Parent: #204 — Blackwell Phase 2 kernel port
Branch: tairan/blackwell-02-04-mxfp4-triton-dispatch
Spec: blackwell-kernel-port-v1.md § Sub-task 4
What
Wire arch == 'sm100' branch in fused_wgmma_grouped.py and fused_wgmma_expert.py to call the existing pure-Triton fused_mxfp4_gemm.py (no WGMMA intrinsics, JIT-compiles for sm100). Tune BLOCK_M/N for B200 128KB SMEM.
Acceptance
- MXFP4 grouped and expert forward: max_abs_err < 5e-3 vs float32 reference
- No _C_expert_mxfp4_wgmma / _C_grouped_mxfp4_wgmma import at runtime on sm100
Parent: #204 — Blackwell Phase 2 kernel port
Branch: tairan/blackwell-02-04-mxfp4-triton-dispatch
Spec: blackwell-kernel-port-v1.md § Sub-task 4
What
Wire arch == 'sm100' branch in fused_wgmma_grouped.py and fused_wgmma_expert.py to call the existing pure-Triton fused_mxfp4_gemm.py (no WGMMA intrinsics, JIT-compiles for sm100). Tune BLOCK_M/N for B200 128KB SMEM.
Acceptance