mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-10-21 06:40:23 +08:00
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.