SIGN IN SIGN UP

[CUDA] Bridge half-style math intrinsics for 16-bit CUTLASS types (#3077)

TVM lowers 16-bit math operations to CUDA's half-style names, but CUDA
provides overloads only for native __half and __nv_bfloat16 values.
Calls with CUTLASS half_t and bfloat16_t wrappers therefore failed to
compile.

Add explicit forwarding overloads for hexp, hexp2, hexp10, hlog, hlog2,
hlog10, hsin, hcos, hfloor, hceil, hrint, and htrunc. Each overload
converts the wrapper to its native CUDA type, calls the existing
intrinsic, and wraps the result back in the original type. Fold the
previous standalone bfloat16 hexp bridge into the complete overload
set. Codegen and the fast-math path remain unchanged.

Add scalar runtime coverage for both wrapper types using materialized
temporaries, separate valid input domains, generated-call assertions,
and torch references.
C
Chennes committed
fe2951323fb79c6e5b909da51333e1d369bb0529
Parent: e8b9da9
Committed by GitHub <noreply@github.com> on 8/25/2026, 6:49:25 AM