From 424fe6bf2c1b9d63350002821025b8a9ff1c15d1 Mon Sep 17 00:00:00 2001
From: biondizzle <biondizzle@gmail.com>
Date: Mon, 1 Jun 2026 07:44:53 +0000
Subject: [PATCH] Fix: use SM100_MMA_MXF8F6F4_SS (not MXF4) to match
 Nvfp4Linear path
MIME-Version: 1.0
Content-Type: text/plain; charset=UTF-8
Content-Transfer-Encoding: 8bit

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.
---
 dsv4/kernels/router/nvfp4_fused_router_kernel.cuh | 7 ++++++-
 1 file changed, 6 insertions(+), 1 deletion(-)

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);