[CUDA] Bridge half-style math intrinsics for 16-bit CUTLASS types

[Chennesxu]

Summary

Many unary T.* math ops fail on float16 and bfloat16 because CUDA lowering emits half-style names such as hlog and hfloor, while CUDA only overloads them for native __half and __nv_bfloat16, not CUTLASS wrappers.

Fix

Add flat half_t and bfloat16_t forwarding overloads for: hexp, hexp2, hexp10, hlog, hlog2, hlog10, hsin, hcos, hfloor, hceil, hrint, and htrunc. Each overload calls the existing native CUDA intrinsic and wraps the result back in the original type. The existing bfloat16 hexp overload is retained as part of the completed set. Codegen and the fast-math path are unchanged.

This enables T.exp, T.exp2, T.exp10, T.log, T.log2, T.log10, T.sin, T.cos, T.floor, T.ceil, T.round, T.trunc, and T.sigmoid through its exp lowering.

Tests

Added one kernel per dtype covering all 12 intrinsics with materialized temporaries, valid positive/mixed-sign inputs, entry-body codegen assertions, and torch references.

Verified locally on sm_75 / CUDA 12.4:

• Touched test file: 5 passed.
• Negative control: 2 failed, 3 passed.
• Related fast-math, clamp, warp-sync, and intrinsic suites: 138 passed.

Scope

Limited to half-style names that already have native CUDA intrinsics. Operations requiring a float32 fallback, missing lowering rules, or separate float8/integer dtype policy remain out of scope.

Summary

• Add CUDA forwarding overloads for 12 half-style math intrinsics on half_t and bfloat16_t.
• Convert wrapper types to native CUDA types, call the existing intrinsic, and convert the result back.
• Extend coverage to exponential, logarithmic, trigonometric, rounding, and truncation operations.
• Add code generation and Torch reference tests for both data types.
• Cover materialized temporaries and verify each lowered CUDA intrinsic.

Validation

• Touched tests and related suites passed locally on sm_75 with CUDA 12.4.
• Code generation and fast-math behavior remain unchanged.

C++ style / lint notes

• The PR changes C++ overloads in src/tl_templates/cuda/common.h.
• It does not change the rules documented in docs/developer_guide/cpp_style.md.
• The “C++ API Style Audit (warning only)” CI step may report advisory findings. Such warnings are not correctness or build failures and should not block this change without a clear API, FFI, or maintainability risk.

[github-actions]

👋 Hi! Thank you for contributing to the TileLang project.

Please remember to run pre-commit run --all-files in the root directory of the project to ensure your changes are properly linted and formatted. This will help ensure your contribution passes the format check.

We appreciate you taking this step! Our team will review your contribution, and we look forward to your awesome work! 🚀

[coderabbitai]

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Pro Plus

Run ID: 04be6c1d-29e1-4c0e-981c-24d53959bd3e

📥 Commits

Reviewing files that changed from the base of the PR and between fded655 and b65abab.

📒 Files selected for processing (2)

• src/tl_templates/cuda/common.h
• testing/python/language/test_tilelang_language_intrinsics_codegen.py

Included review availability: Your plan provides up to 8 included reviews per hour; 7 remain after this review.

---

📝 Walkthrough

Walkthrough

Changes

CUDA 16-bit math support

Layer / File(s)	Summary
CUDA intrinsic bridges src/tl_templates/cuda/common.h	Adds half_t and bfloat16_t wrappers for exponential, logarithmic, trigonometric, and rounding CUDA intrinsics.
Kernel code-generation validation testing/python/language/test_tilelang_language_intrinsics_codegen.py	Adds metadata-driven kernels and CUDA tests for generated intrinsic names and PyTorch output comparisons.

Estimated code review effort: 3 (Moderate) | ~20 minutes

Merge Risk: ⚪ Minimal · up to b65ab

This localized change adds half-style math forwarding for CUTLASS 16-bit types and includes focused coverage; no actionable merge-blocking risk remains beyond normal checks and review.

Sequence Diagram(s)

sequenceDiagram
  participant Test
  participant TileLangCompiler
  participant CUDA
  participant PyTorch
  Test->>TileLangCompiler: Compile half_math_kernel
  TileLangCompiler-->>Test: Return generated CUDA
  Test->>CUDA: Execute kernel
  CUDA-->>Test: Return outputs
  Test->>PyTorch: Compute reference outputs
  PyTorch-->>Test: Return references

Loading

Suggested reviewers: leiwang1999, rachmanino

🚥 Pre-merge checks | ✅ 5

✅ Passed checks (5 passed)

Check name	Status	Explanation
Description Check	✅ Passed	Check skipped - CodeRabbit’s high-level summary is enabled.
Title check	✅ Passed	The title clearly summarizes the main change: adding CUDA bridges for half-style math intrinsics on 16-bit CUTLASS types.
Docstring Coverage	✅ Passed	Docstring check was indeterminate for this PR — some files could not be analyzed in time. Not blocking.
Linked Issues check	✅ Passed	Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check	✅ Passed	Check skipped because no linked issues were found for this pull request.

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create PR with unit tests

---

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

• X
• Mastodon
• Reddit
• LinkedIn

_{Comment @coderabbitai help to get the list of available commands.}