diff --git a/dsv4/kernels/router/nvfp4_fused_router_kernel.cuh b/dsv4/kernels/router/nvfp4_fused_router_kernel.cuh index db9148b3..15dd90d4 100644 --- a/dsv4/kernels/router/nvfp4_fused_router_kernel.cuh +++ b/dsv4/kernels/router/nvfp4_fused_router_kernel.cuh @@ -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);