mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-10-21 06:40:23 +08:00
Chennes
7fd95363cc
On SM75, SelectInst routed GEMMs to the plain MMA path even when their dtype combination had no native Turing mma.sync atom. In particular, fp32, fp64, and bf16 selected SM80-only atoms: bf16 failed an nvcc static_assert, while fp32 and fp64 compiled into runtime traps that invalidated the CUDA context. Add AllowTuringMma alongside the Volta check introduced in #2339 and select the existing cuda.fma implementation when SM75 has no native atom. Keep the check dtype-only because the SM75 emitter supports the same operand scopes and transposes as the generic MMA path. Valid f16 combinations remain on native MMA rather than being demoted to the slower and less accurate serial FMA fallback. Add compile-only dispatch checks pinned to an explicit sm_75 target, so they run on any CUDA runner, verifying that fp32, fp64, and bf16 use the fallback while f16 remains on MMA. Add execution checks gated on compute capability 7.5.