Files
tilelang/testing/python/math
Chennes c6294f07e3 [BugFix] Add pre-SM80 fallback for bf16 __hfma (#2769)
common.h defines a bfloat16_t __hfma wrapper that forwards to the
__nv_bfloat16 __hfma overload. CUDA only declares that overload for
__CUDA_ARCH__ >= 800, so on SM75 merely parsing the wrapper fails with
"no instance of overloaded function __hfma". Because common.h is included
by every generated kernel, this breaks all SM75 builds, not only bf16 ones.

Guard the native path on SM80+ and use an fp32 FMA followed by bf16
conversion on earlier architectures, following the existing pre-SM80 bf16
fallback strategy.

Add a compile-only regression test that builds a bf16 ieee_fmaf kernel for
an explicit SM75 target. Existing IEEE math tests use target="cuda", so an
SM80+ runner exercises only the native branch.
2026-07-25 19:57:47 +08:00
..