SIGN IN SIGN UP

[CUDA] Add 16-bit overloads for CUTLASS fast-math functions (#3097)

math.h aliases several half-style names to cutlass::fast_* when a kernel
uses selected fast-math intrinsics. CUTLASS provides only float and double
overloads for fast_log, fast_sqrt, fast_sin, and fast_cos, and adds half_t
overloads only for fast_exp and fast_tanh. Overload resolution can therefore
leave a float result that cannot copy-initialize the CUTLASS wrapper.

Once a kernel includes math.h, for example through T.__exp, these aliases
also affect matching plain math calls in the same kernel. Their compilation
could therefore depend on whether another intrinsic pulled math.h in.

Add half_t and bfloat16_t overloads for fast_log, fast_sqrt, fast_sin, and
fast_cos, plus bfloat16_t for fast_tanh. Each evaluates in float32 and
converts back, following the existing fast_exp(bfloat16_t) bridge. Calls
that already compiled retain the same numerical path.

Add compile-and-run coverage for both wrapper types.
C
Chennes committed
035a2af35b4d59f12fa1d2572b2eef821b62923a
Parent: 966e646
Committed by GitHub <noreply@github.com> on 8/27/2026, 9:17:24 AM