mirror of
https://github.com/tile-ai/tilelang.git
synced 2026-11-01 08:23:24 +08:00
Merging the SM120 block-scale mode into TensorCoreIntrinEmitter changed mma()'s 4th/5th positional parameters to SFA_buf/SFB_buf, so the plain gemm lowering's positional mma(A_local, B_local, C_buf, ki) call bound ki to SFA_buf and raised 'Scale buffers require TensorCoreIntrinEmitter block-scale mode' for every non-blockscaled MMA GEMM (CI examples: attention_sink, deepseek_nsa/mhc/v32). Keep the base (A, B, C, k_inner) positional signature, make all scale-related parameters keyword-only, forward k_inner to the base implementation, and drop the unused blockscaled mma_atom override that shadowed the base method the same way. Add a signature-contract regression test. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>