biondizzle/nvfp4-megamoe-kernel
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Fix: use SM100_MMA_MXF8F6F4_SS (not MXF4) to match Nvfp4Linear path

Browse Source 
MXF4 has .block32 hardcoded. MXF8F6F4 matches what CuTeDSL generates
via make_instr_desc_block_scaled. Both use E2M1 data + UE8M0 scales
at hardware level. NVFP4 E2M1 microscales are combined into UE8M0
during quantization — no MXFP4 conversion.
This commit is contained in:
biondizzle
2026-06-01 07:44:53 +00:00
parent 2e2caadf7d
commit 424fe6bf2c
1 changed files with 6 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
dsv4/kernels/router/nvfp4_fused_router_kernel.cuh
View File

@@ -393,7 +393,12 @@ void nvfp4_fused_router_kernel(
cute::UMMA::Major::K, BLOCK_N, kSwizzleBMode, b_dtype_t>(
b_desc_base_lo, 0, k * UMMA_K);
deep_gemm::ptx::SM100_MMA_MXF4_SS::fma(
// Use MXF8F6F4 (not MXF4) to match Nvfp4Linear's
// make_instr_desc_block_scaled path. Both use E2M1 data
// + UE8M0 scales at the hardware level. The NVFP4
// E2M1 microscales are combined into UE8M0 during
// quantization — no MXFP4 conversion needed.
deep_gemm::ptx::SM100_MMA_MXF8F6F4_SS::fma(
a_desc, b_desc, accum_stage * UMMA_N,
first_k ? 0 : 1, runtime_desc,
kTmemStartColOfSFA, kTmemStartColOfSFB);