diff --git a/cpp/tensorrt_llm/kernels/mhcKernels/fused_tf32_pmap_gemm.cuh b/cpp/tensorrt_llm/kernels/mhcKernels/fused_tf32_pmap_gemm.cuh
index 9e74d4eb20f7..bdbdd1a44994 100644
--- a/cpp/tensorrt_llm/kernels/mhcKernels/fused_tf32_pmap_gemm.cuh
+++ b/cpp/tensorrt_llm/kernels/mhcKernels/fused_tf32_pmap_gemm.cuh
@@ -108,11 +108,11 @@ __device__ __forceinline__ void stsm_x4_b16_rout(void* smem_dst, uint32_t a, uin
 
 template <uint32_t SHAPE_N, uint32_t HIDDEN, uint32_t HC_MULT, uint32_t BLOCK_M, uint32_t BLOCK_N, uint32_t BLOCK_K,
     uint32_t kSwizzleCDMode, uint32_t N_B_STAGES, uint32_t N_INPUT_STAGES, uint32_t kNumMMAThreads,
-    uint32_t kNumPmapThreads, uint32_t kNumSplits = 1, bool kEarlyRelease = false>
+    uint32_t kNumPmapThreads, uint32_t kNumSplits = 1, bool kEarlyRelease = false, uint32_t kXSplit = 1>
 __global__ void __launch_bounds__(kNumMMAThreads + kNumPmapThreads, 1) fused_tf32_pmap_gemm_rout_atomic_impl(
-    const uint32_t shape_m, const __grid_constant__ cute::TmaDescriptor tensor_map_residual,
-    const __grid_constant__ cute::TmaDescriptor tensor_map_x, const __grid_constant__ cute::TmaDescriptor tensor_map_b,
-    const __grid_constant__ cute::TmaDescriptor tensor_map_residual_out,
+    uint32_t const shape_m, __grid_constant__ const cute::TmaDescriptor tensor_map_residual,
+    __grid_constant__ const cute::TmaDescriptor tensor_map_x, __grid_constant__ const cute::TmaDescriptor tensor_map_b,
+    __grid_constant__ const cute::TmaDescriptor tensor_map_residual_out,
     float* __restrict__ D, // [M, SHAPE_N]  (caller memsets to 0)
     float const* __restrict__ post_mix, float const* __restrict__ comb_mix, float* __restrict__ sqr_sum)
 {                          // [M]            (caller memsets to 0)
@@ -175,6 +175,7 @@ __global__ void __launch_bounds__(kNumMMAThreads + kNumPmapThreads, 1) fused_tf3
     auto smem_res = PatternVisitor(
         [&, base = cursor](uint32_t const& i) { return reinterpret_cast<nv_bfloat16*>(base + i * SMEM_RES_PER_ISTG); });
     cursor += N_INPUT_STAGES * SMEM_RES_PER_ISTG;
+    // Stream split-x tiles through the existing x buffers.
     auto smem_x_stg = PatternVisitor(
         [&, base = cursor](uint32_t const& i) { return reinterpret_cast<nv_bfloat16*>(base + i * SMEM_X_PER_ISTG); });
     cursor += N_INPUT_STAGES * SMEM_X_PER_ISTG;
@@ -192,13 +193,18 @@ __global__ void __launch_bounds__(kNumMMAThreads + kNumPmapThreads, 1) fused_tf3
     auto full_input = PatternVisitor([=](uint32_t const& i) { return barrier_start_ptr + 2 * N_B_STAGES + i; });
     auto empty_input
         = PatternVisitor([=](uint32_t const& i) { return barrier_start_ptr + 2 * N_B_STAGES + N_INPUT_STAGES + i; });
-    auto full_cast = PatternVisitor(
+    constexpr uint32_t kXBarriers = (kXSplit > 1) ? 2 * N_INPUT_STAGES : 0;
+    auto full_x = PatternVisitor(
         [=](uint32_t const& i) { return barrier_start_ptr + 2 * N_B_STAGES + 2 * N_INPUT_STAGES + i; });
+    auto empty_x = PatternVisitor(
+        [=](uint32_t const& i) { return barrier_start_ptr + 2 * N_B_STAGES + 3 * N_INPUT_STAGES + i; });
+    auto full_cast = PatternVisitor(
+        [=](uint32_t const& i) { return barrier_start_ptr + 2 * N_B_STAGES + 2 * N_INPUT_STAGES + kXBarriers + i; });
     auto empty_cast = PatternVisitor([=](uint32_t const& i)
-        { return barrier_start_ptr + 2 * N_B_STAGES + 2 * N_INPUT_STAGES + kNumCastStages + i; });
-    auto tmem_full_barrier = barrier_start_ptr + 2 * N_B_STAGES + 2 * N_INPUT_STAGES + 2 * kNumCastStages;
+        { return barrier_start_ptr + 2 * N_B_STAGES + 2 * N_INPUT_STAGES + kXBarriers + kNumCastStages + i; });
+    auto tmem_full_barrier = barrier_start_ptr + 2 * N_B_STAGES + 2 * N_INPUT_STAGES + kXBarriers + 2 * kNumCastStages;
 
-    cursor += (2 * N_B_STAGES + 2 * N_INPUT_STAGES + 2 * kNumCastStages + 1) * sizeof(Barrier);
+    cursor += (2 * N_B_STAGES + 2 * N_INPUT_STAGES + kXBarriers + 2 * kNumCastStages + 1) * sizeof(Barrier);
     auto tmem_ptr_in_smem = reinterpret_cast<uint32_t*>(cursor);
 
     if (warp_idx == 1 and cute::elect_one_sync())
@@ -215,6 +221,15 @@ __global__ void __launch_bounds__(kNumMMAThreads + kNumPmapThreads, 1) fused_tf3
             full_input[i]->init(1);
             empty_input[i]->init(kNumPmapThreads);
         }
+        if constexpr (kXSplit > 1)
+        {
+#pragma unroll
+            for (uint32_t i = 0; i < N_INPUT_STAGES; ++i)
+            {
+                full_x[i]->init(1);
+                empty_x[i]->init(kNumPmapThreads);
+            }
+        }
 #pragma unroll
         for (uint32_t i = 0; i < kNumCastStages; ++i)
         {
@@ -278,6 +293,7 @@ __global__ void __launch_bounds__(kNumMMAThreads + kNumPmapThreads, 1) fused_tf3
             uint32_t b_stage = 0;
             uint32_t i_stage = 0;
             uint32_t s = 0;
+            uint32_t xs = 0;
             for (uint32_t ht = 0; ht < H_TILES_PER_SPLIT; ++ht)
             {
                 const uint32_t h_tile = h_tile_start + ht;
@@ -290,10 +306,27 @@ __global__ void __launch_bounds__(kNumMMAThreads + kNumPmapThreads, 1) fused_tf3
                     tma_copy<BLOCK_K, BLOCK_M, kSwizzleResMode>(&tensor_map_residual, full_input[i_stage],
                         smem_res[i_stage] + j * BLOCK_M * BLOCK_K, j * HIDDEN + h_idx, m_idx);
                 }
-                tma_copy<BLOCK_K, BLOCK_M, kSwizzleXMode>(
-                    &tensor_map_x, full_input[i_stage], smem_x_stg[i_stage], h_idx, m_idx);
-                constexpr uint32_t kInputBytes = SMEM_RES_PER_ISTG + SMEM_X_PER_ISTG;
-                full_input[i_stage]->arrive_and_expect_tx(kInputBytes);
+                if constexpr (kXSplit == 1)
+                {
+                    tma_copy<BLOCK_K, BLOCK_M, kSwizzleXMode>(
+                        &tensor_map_x, full_input[i_stage], smem_x_stg[i_stage], h_idx, m_idx);
+                    constexpr uint32_t kInputBytes = SMEM_RES_PER_ISTG + SMEM_X_PER_ISTG;
+                    full_input[i_stage]->arrive_and_expect_tx(kInputBytes);
+                }
+                else
+                {
+                    full_input[i_stage]->arrive_and_expect_tx(SMEM_RES_PER_ISTG);
+#pragma unroll
+                    for (uint32_t xp = 0; xp < kXSplit; ++xp)
+                    {
+                        uint32_t const xslot = xs % N_INPUT_STAGES;
+                        empty_x[xslot]->wait(((xs / N_INPUT_STAGES) & 1) ^ 1);
+                        tma_copy<BLOCK_K, BLOCK_M, kSwizzleXMode>(
+                            &tensor_map_x, full_x[xslot], smem_x_stg[xslot], h_idx, xp * shape_m + m_idx);
+                        full_x[xslot]->arrive_and_expect_tx(SMEM_X_PER_ISTG);
+                        ++xs;
+                    }
+                }
 
 #pragma unroll
                 for (uint32_t hc = 0; hc < HC_MULT; ++hc)
@@ -448,19 +481,6 @@ __global__ void __launch_bounds__(kNumMMAThreads + kNumPmapThreads, 1) fused_tf3
             const uint32_t i_stage = ht % N_INPUT_STAGES;
             full_input[i_stage]->wait((ht / N_INPUT_STAGES) & 1);
 
-            uint32_t x_vals[2][kNumLoads];
-            {
-                uint8_t const* x_base
-                    = reinterpret_cast<uint8_t*>(smem_x_stg[i_stage]) + sub_warp_idx * BLOCK_M_PER_WARP * kSwizzleXMode;
-#pragma unroll
-                for (uint32_t i = 0; i < kNumLoads; i += 2)
-                {
-                    auto smem_ptr = x_base + get_swizzled_smem_offset<kSwizzleXMode>(i + lane_idx / 16, lane_idx % 16);
-                    deep_gemm::sm90::SM90_U32x4_LDSM_N::copy(x_vals[0][i + 0], x_vals[1][i + 0], x_vals[0][i + 1],
-                        x_vals[1][i + 1], const_cast<uint8_t*>(smem_ptr));
-                }
-            }
-
             uint32_t r_vals[HC_MULT][2][kNumLoads];
 #pragma unroll
             for (uint32_t j = 0; j < HC_MULT; ++j)
@@ -480,12 +500,48 @@ __global__ void __launch_bounds__(kNumMMAThreads + kNumPmapThreads, 1) fused_tf3
             float2 xf[2][kNumLoads];
 #pragma unroll
             for (uint32_t u = 0; u < 2; ++u)
-            {
 #pragma unroll
                 for (uint32_t i = 0; i < kNumLoads; ++i)
+                    xf[u][i] = make_float2(0.f, 0.f);
+#pragma unroll
+            for (uint32_t xp = 0; xp < kXSplit; ++xp)
+            {
+                nv_bfloat16* x_buf;
+                if constexpr (kXSplit == 1)
                 {
-                    xf[u][i] = __bfloat1622float2(*reinterpret_cast<nv_bfloat162*>(&x_vals[u][i]));
+                    x_buf = smem_x_stg[i_stage];
+                }
+                else
+                {
+                    uint32_t const xsg = ht * kXSplit + xp;
+                    uint32_t const xslot = xsg % N_INPUT_STAGES;
+                    full_x[xslot]->wait((xsg / N_INPUT_STAGES) & 1);
+                    x_buf = smem_x_stg[xslot];
+                }
+
+                uint32_t x_vals[2][kNumLoads];
+                uint8_t const* x_base
+                    = reinterpret_cast<uint8_t*>(x_buf) + sub_warp_idx * BLOCK_M_PER_WARP * kSwizzleXMode;
+#pragma unroll
+                for (uint32_t i = 0; i < kNumLoads; i += 2)
+                {
+                    auto smem_ptr = x_base + get_swizzled_smem_offset<kSwizzleXMode>(i + lane_idx / 16, lane_idx % 16);
+                    deep_gemm::sm90::SM90_U32x4_LDSM_N::copy(x_vals[0][i + 0], x_vals[1][i + 0], x_vals[0][i + 1],
+                        x_vals[1][i + 1], const_cast<uint8_t*>(smem_ptr));
+                }
+                if constexpr (kXSplit > 1)
+                {
+                    empty_x[(ht * kXSplit + xp) % N_INPUT_STAGES]->arrive();
                 }
+#pragma unroll
+                for (uint32_t u = 0; u < 2; ++u)
+#pragma unroll
+                    for (uint32_t i = 0; i < kNumLoads; ++i)
+                    {
+                        float2 value = __bfloat1622float2(*reinterpret_cast<nv_bfloat162*>(&x_vals[u][i]));
+                        xf[u][i].x += value.x;
+                        xf[u][i].y += value.y;
+                    }
             }
 
             if constexpr (kEarlyRelease)
diff --git a/cpp/tensorrt_llm/kernels/mhcKernels/mhcFusedHcKernel.cu b/cpp/tensorrt_llm/kernels/mhcKernels/mhcFusedHcKernel.cu
index 86586edc099a..a38db0e55752 100644
--- a/cpp/tensorrt_llm/kernels/mhcKernels/mhcFusedHcKernel.cu
+++ b/cpp/tensorrt_llm/kernels/mhcKernels/mhcFusedHcKernel.cu
@@ -293,14 +293,53 @@ static constexpr uint32_t fhcSmemSize()
 using FusedRoutFn = void (*)(
     uint32_t, CUtensorMap, CUtensorMap, CUtensorMap, CUtensorMap, float*, float const*, float const*, float*);
 
-template <uint32_t Hidden, uint32_t KS>
+template <uint32_t Hidden, uint32_t KS, uint32_t XS = 1>
 static FusedRoutFn fhcInstance()
 {
     static_assert(isSupportedFhcHidden<Hidden>(), "Unsupported fused-HC hidden size");
     static_assert(isSupportedFhcMmaKS<Hidden, KS>(), "Unsupported fused-HC MMA kNumSplits for hidden size");
     return &fused_mhc::fused_tf32_pmap_gemm_rout_atomic_impl<FHC_SHAPE_N, Hidden, FHC_HC_MULT, FHC_BLOCK_M, FHC_BLOCK_N,
         FHC_BLOCK_K, FHC_SWIZZLE_CD, FHC_N_B_STAGES, FHC_N_INPUT_STG, FHC_NUM_MMA_TH, FHC_NUM_PMAP_TH, KS,
-        /*kEarlyRelease=*/false>;
+        /*kEarlyRelease=*/false, XS>;
+}
+
+template <uint32_t Hidden, uint32_t KS, uint32_t XS>
+static FusedRoutFn fhcXSplitInstanceIfSupported()
+{
+    if constexpr (isSupportedFhcMmaKS<Hidden, KS>())
+    {
+        return fhcInstance<Hidden, KS, XS>();
+    }
+    else
+    {
+        TLLM_CHECK_WITH_INFO(false, "mhcFusedHcLaunch: unsupported (kNumSplits=%u, hidden=%u)", KS, Hidden);
+        return nullptr;
+    }
+}
+
+template <uint32_t Hidden, uint32_t XS>
+static FusedRoutFn pickFhcXSplitKs(uint32_t ks)
+{
+    switch (ks)
+    {
+    case 1: return fhcInstance<Hidden, 1, XS>();
+    case 2: return fhcXSplitInstanceIfSupported<Hidden, 2, XS>();
+    case 4: return fhcXSplitInstanceIfSupported<Hidden, 4, XS>();
+    case 8: return fhcXSplitInstanceIfSupported<Hidden, 8, XS>();
+    case 16: return fhcXSplitInstanceIfSupported<Hidden, 16, XS>();
+    default: TLLM_CHECK_WITH_INFO(false, "mhcFusedHcLaunch: unsupported kNumSplits=%u for split x", ks); return nullptr;
+    }
+}
+
+template <uint32_t Hidden>
+static FusedRoutFn pickFhcXSplit(uint32_t ks, uint32_t xs)
+{
+    switch (xs)
+    {
+    case 2: return pickFhcXSplitKs<Hidden, 2>(ks);
+    case 4: return pickFhcXSplitKs<Hidden, 4>(ks);
+    default: TLLM_CHECK_WITH_INFO(false, "mhcFusedHcLaunch: unsupported x split=%u", xs); return nullptr;
+    }
 }
 
 template <uint32_t Hidden, uint32_t KS>
@@ -375,7 +414,7 @@ static void mhcFusedHcLaunchImpl(__nv_bfloat16 const* x_prev, __nv_bfloat16 cons
     __nv_bfloat16* layer_input_cur, float* y_acc_workspace, float* r_acc_workspace, int M, int hidden_size, int hc_mult,
     int num_k_splits, int bigfuse_block_size, float rms_eps, float hc_pre_eps, float hc_sinkhorn_eps,
     float hc_post_mult_value, int sinkhorn_repeat, __nv_bfloat16 const* norm_weight, float norm_eps,
-    cudaStream_t stream)
+    cudaStream_t stream, int x_num_splits = 1)
 {
     if (M <= 0)
         return;
@@ -407,7 +446,8 @@ static void mhcFusedHcLaunchImpl(__nv_bfloat16 const* x_prev, __nv_bfloat16 cons
         /*swizzleBytes=*/128, sizeof(__nv_bfloat16));
 
     CUtensorMap desc_x = getCachedTma2D(const_cast<__nv_bfloat16*>(x_prev), CU_TENSOR_MAP_DATA_TYPE_BFLOAT16, Hidden,
-        m_u, FHC_BLOCK_K, FHC_BLOCK_M, static_cast<uint64_t>(Hidden) * sizeof(__nv_bfloat16),
+        static_cast<uint32_t>(x_num_splits) * m_u, FHC_BLOCK_K, FHC_BLOCK_M,
+        static_cast<uint64_t>(Hidden) * sizeof(__nv_bfloat16),
         /*swizzleBytes=*/128, sizeof(__nv_bfloat16));
 
     CUtensorMap desc_b = getCachedTma2D(const_cast<float*>(w_t), CU_TENSOR_MAP_DATA_TYPE_TFLOAT32, SHAPE_K, FHC_SHAPE_N,
@@ -419,8 +459,11 @@ static void mhcFusedHcLaunchImpl(__nv_bfloat16 const* x_prev, __nv_bfloat16 cons
         /*swizzleBytes=*/128, sizeof(__nv_bfloat16));
 
     // ---- Step 1: fused post-mapping + TF32 GEMM + sqrsum + residual_out ----
-    constexpr uint32_t fused_smem = fhcSmemSize();
-    FusedRoutFn fa = pickFhc<Hidden>(ks);
+    // Split-x adds barriers but reuses the x data buffers.
+    uint32_t const extra_x_smem = (x_num_splits > 1) ? (2u * FHC_N_INPUT_STG * sizeof(uint64_t)) : 0u;
+    uint32_t const fused_smem = fhcSmemSize() + extra_x_smem;
+    FusedRoutFn fa
+        = (x_num_splits > 1) ? pickFhcXSplit<Hidden>(ks, static_cast<uint32_t>(x_num_splits)) : pickFhc<Hidden>(ks);
     TLLM_CUDA_CHECK(cudaFuncSetAttribute(
         reinterpret_cast<void const*>(fa), cudaFuncAttributeMaxDynamicSharedMemorySize, fused_smem));
 
@@ -443,7 +486,7 @@ void mhcFusedHcLaunch(__nv_bfloat16 const* x_prev, __nv_bfloat16 const* residual
     __nv_bfloat16* residual_cur, float* post_mix_cur, float* comb_mix_cur, __nv_bfloat16* layer_input_cur,
     float* y_acc_workspace, float* r_acc_workspace, int M, int hidden_size, int hc_mult, int num_k_splits,
     int bigfuse_block_size, float rms_eps, float hc_pre_eps, float hc_sinkhorn_eps, float hc_post_mult_value,
-    int sinkhorn_repeat, __nv_bfloat16 const* norm_weight, float norm_eps, cudaStream_t stream)
+    int sinkhorn_repeat, __nv_bfloat16 const* norm_weight, float norm_eps, cudaStream_t stream, int x_num_splits)
 {
     if (M <= 0)
         return;
@@ -457,13 +500,13 @@ void mhcFusedHcLaunch(__nv_bfloat16 const* x_prev, __nv_bfloat16 const* residual
         mhcFusedHcLaunchImpl<FHC_HIDDEN_FLASH>(x_prev, residual_prev, post_mix_prev, comb_mix_prev, w_t, hc_scale,
             hc_base, residual_cur, post_mix_cur, comb_mix_cur, layer_input_cur, y_acc_workspace, r_acc_workspace, M,
             hidden_size, hc_mult, num_k_splits, bigfuse_block_size, rms_eps, hc_pre_eps, hc_sinkhorn_eps,
-            hc_post_mult_value, sinkhorn_repeat, norm_weight, norm_eps, stream);
+            hc_post_mult_value, sinkhorn_repeat, norm_weight, norm_eps, stream, x_num_splits);
         return;
     case static_cast<int>(FHC_HIDDEN_PRO):
         mhcFusedHcLaunchImpl<FHC_HIDDEN_PRO>(x_prev, residual_prev, post_mix_prev, comb_mix_prev, w_t, hc_scale,
             hc_base, residual_cur, post_mix_cur, comb_mix_cur, layer_input_cur, y_acc_workspace, r_acc_workspace, M,
             hidden_size, hc_mult, num_k_splits, bigfuse_block_size, rms_eps, hc_pre_eps, hc_sinkhorn_eps,
-            hc_post_mult_value, sinkhorn_repeat, norm_weight, norm_eps, stream);
+            hc_post_mult_value, sinkhorn_repeat, norm_weight, norm_eps, stream, x_num_splits);
         return;
     default: return;
     }
@@ -481,19 +524,21 @@ void mhcFusedHcLaunch(__nv_bfloat16 const* x_prev, __nv_bfloat16 const* residual
 using FmaKsplitFn = void (*)(__nv_bfloat16 const*, __nv_bfloat16 const*, float const*, float const*, float const*,
     float*, float*, int, int, int, __nv_bfloat16*);
 
-template <int TN, int KS>
+template <int TN, int KS, int XS = 1>
 static FmaKsplitFn fhcFmaInstance()
 {
-    return &fused_fma_kernels::fused_pmap_gemm_fma_ksplit<TN, KS, /*BF16_VEC_OVERRIDE=*/0, /*WRITE_RESIDUAL=*/true>;
+    return &fused_fma_kernels::fused_pmap_gemm_fma_ksplit<TN, KS, /*BF16_VEC_OVERRIDE=*/0,
+        /*WRITE_RESIDUAL=*/true, XS>;
 }
 
 // Valid (tile_n, num_k_splits) combinations the fused_hc FMA path supports.
 // Keep this limited to the small/mid-M sweet spots from profile_fair_report v4.
-static FmaKsplitFn pickFhcFma(int tile_n, int ks)
+template <int XS>
+static FmaKsplitFn pickFhcFmaXSplit(int tile_n, int ks)
 {
 #define FHCFMA_CASE(TN, KS)                                                                                            \
     if (tile_n == (TN) && ks == (KS))                                                                                  \
-    return fhcFmaInstance<TN, KS>()
+    return fhcFmaInstance<TN, KS, XS>()
 
     FHCFMA_CASE(1, 1);
     FHCFMA_CASE(1, 2);
@@ -513,16 +558,29 @@ static FmaKsplitFn pickFhcFma(int tile_n, int ks)
     FHCFMA_CASE(12, 1);
     FHCFMA_CASE(24, 1);
 #undef FHCFMA_CASE
-    TLLM_CHECK_WITH_INFO(false, "mhcFusedHcFmaLaunch: unsupported (tile_n=%d, ks=%d)", tile_n, ks);
+    TLLM_CHECK_WITH_INFO(false, "mhcFusedHcFmaLaunch: unsupported (tile_n=%d, ks=%d, x_splits=%d)", tile_n, ks, XS);
     return nullptr;
 }
 
+static FmaKsplitFn pickFhcFma(int tile_n, int ks, int x_num_splits)
+{
+    switch (x_num_splits)
+    {
+    case 1: return pickFhcFmaXSplit<1>(tile_n, ks);
+    case 2: return pickFhcFmaXSplit<2>(tile_n, ks);
+    case 4: return pickFhcFmaXSplit<4>(tile_n, ks);
+    default:
+        TLLM_CHECK_WITH_INFO(false, "mhcFusedHcFmaLaunch: unsupported x_num_splits=%d", x_num_splits);
+        return nullptr;
+    }
+}
+
 void mhcFusedHcFmaLaunch(__nv_bfloat16 const* x_prev, __nv_bfloat16 const* residual_prev, float const* post_mix_prev,
     float const* comb_mix_prev, float const* w_t, float const* hc_scale, float const* hc_base,
     __nv_bfloat16* residual_cur, float* post_mix_cur, float* comb_mix_cur, __nv_bfloat16* layer_input_cur,
     float* y_acc_workspace, float* r_acc_workspace, int M, int hidden_size, int hc_mult, int tile_n, int num_k_splits,
     int bigfuse_block_size, float rms_eps, float hc_pre_eps, float hc_sinkhorn_eps, float hc_post_mult_value,
-    int sinkhorn_repeat, __nv_bfloat16 const* norm_weight, float norm_eps, cudaStream_t stream)
+    int sinkhorn_repeat, __nv_bfloat16 const* norm_weight, float norm_eps, cudaStream_t stream, int x_num_splits)
 {
     if (M <= 0)
         return;
@@ -538,7 +596,7 @@ void mhcFusedHcFmaLaunch(__nv_bfloat16 const* x_prev, __nv_bfloat16 const* resid
     int const N = static_cast<int>(FHC_SHAPE_N);
 
     // ---- Step 1: fused pmap + GEMM + sqrsum + residual_cur (FMA ksplit) ----
-    FmaKsplitFn fn = pickFhcFma(tile_n, num_k_splits);
+    FmaKsplitFn fn = pickFhcFma(tile_n, num_k_splits, x_num_splits);
     dim3 const grid(static_cast<unsigned>(M), static_cast<unsigned>(N / tile_n), static_cast<unsigned>(num_k_splits));
     dim3 const block(256);
     tensorrt_llm::common::launchWithPdlWhenEnabled("fused_pmap_gemm_fma_ksplit", fn, grid, block, 0, stream,
diff --git a/cpp/tensorrt_llm/kernels/mhcKernels/mhcKernels.h b/cpp/tensorrt_llm/kernels/mhcKernels/mhcKernels.h
index 1240f664d522..bb7b03336997 100644
--- a/cpp/tensorrt_llm/kernels/mhcKernels/mhcKernels.h
+++ b/cpp/tensorrt_llm/kernels/mhcKernels/mhcKernels.h
@@ -18,6 +18,8 @@
 #include <cuda_bf16.h>
 #include <cuda_runtime.h>
 
+#include <cstdint>
+
 #include "tensorrt_llm/common/config.h"
 
 TRTLLM_NAMESPACE_BEGIN
@@ -68,7 +70,7 @@ void mhcFusedHcLaunch(__nv_bfloat16 const* x_prev, __nv_bfloat16 const* residual
     __nv_bfloat16* residual_cur, float* post_mix_cur, float* comb_mix_cur, __nv_bfloat16* layer_input_cur,
     float* y_acc_workspace, float* r_acc_workspace, int M, int hidden_size, int hc_mult, int num_k_splits,
     int bigfuse_block_size, float rms_eps, float hc_pre_eps, float hc_sinkhorn_eps, float hc_post_mult_value,
-    int sinkhorn_repeat, __nv_bfloat16 const* norm_weight, float norm_eps, cudaStream_t stream);
+    int sinkhorn_repeat, __nv_bfloat16 const* norm_weight, float norm_eps, cudaStream_t stream, int x_num_splits = 1);
 
 // FMA-path fused hyper-connection boundary launcher.
 //
@@ -96,7 +98,7 @@ void mhcFusedHcFmaLaunch(__nv_bfloat16 const* x_prev, __nv_bfloat16 const* resid
     __nv_bfloat16* residual_cur, float* post_mix_cur, float* comb_mix_cur, __nv_bfloat16* layer_input_cur,
     float* y_acc_workspace, float* r_acc_workspace, int M, int hidden_size, int hc_mult, int tile_n, int num_k_splits,
     int bigfuse_block_size, float rms_eps, float hc_pre_eps, float hc_sinkhorn_eps, float hc_post_mult_value,
-    int sinkhorn_repeat, __nv_bfloat16 const* norm_weight, float norm_eps, cudaStream_t stream);
+    int sinkhorn_repeat, __nv_bfloat16 const* norm_weight, float norm_eps, cudaStream_t stream, int x_num_splits = 1);
 
 // Single-kernel all-in-one fused hyper-connection boundary launcher (TF32 tcgen05 path).
 //
diff --git a/cpp/tensorrt_llm/kernels/mhcKernels/mhc_fused_fma.cuh b/cpp/tensorrt_llm/kernels/mhcKernels/mhc_fused_fma.cuh
index d823f67ba177..939ac95b11c3 100644
--- a/cpp/tensorrt_llm/kernels/mhcKernels/mhc_fused_fma.cuh
+++ b/cpp/tensorrt_llm/kernels/mhcKernels/mhc_fused_fma.cuh
@@ -50,7 +50,8 @@ namespace fused_fma_kernels
 // the post-mapped residual (bf16 [M, HC_MULT, hidden]) into residual_out, using
 // the same disjoint h-slice it already owns.  When false the residual_out arg
 // may be a nullptr; the parameter stays in the signature for simplicity.
-template <int N_PER_BLOCK, int NUM_K_SPLITS, int BF16_VEC_OVERRIDE = 0, bool WRITE_RESIDUAL = false>
+template <int N_PER_BLOCK, int NUM_K_SPLITS, int BF16_VEC_OVERRIDE = 0, bool WRITE_RESIDUAL = false,
+    int X_NUM_SPLITS = 1>
 __launch_bounds__(256) __global__ void fused_pmap_gemm_fma_ksplit(__nv_bfloat16 const* __restrict__ residual_in,
     __nv_bfloat16 const* __restrict__ x_in, float const* __restrict__ post_mix, float const* __restrict__ comb_mix,
     float const* __restrict__ W_T, float* __restrict__ Yp, float* __restrict__ Rp, int hidden_size, int N, int K,
@@ -58,6 +59,7 @@ __launch_bounds__(256) __global__ void fused_pmap_gemm_fma_ksplit(__nv_bfloat16
 {
     static_assert(NUM_K_SPLITS == 1 || NUM_K_SPLITS == 2 || NUM_K_SPLITS == 4 || NUM_K_SPLITS == 8,
         "ksplit ∈ {1, 2, 4, 8} supported");
+    static_assert(X_NUM_SPLITS == 1 || X_NUM_SPLITS == 2 || X_NUM_SPLITS == 4, "x split ∈ {1, 2, 4} supported");
     constexpr int HC_MULT = 4;
     constexpr int BLOCK_SIZE = 256;
     constexpr int WARP_SIZE = 32;
@@ -111,45 +113,49 @@ __launch_bounds__(256) __global__ void fused_pmap_gemm_fma_ksplit(__nv_bfloat16
     float sqr = 0.0f;
 
     long long const tok_res = static_cast<long long>(token) * HC_MULT * hidden_size;
-    long long const tok_x = static_cast<long long>(token) * hidden_size;
-
     // Iterate h over this CTA's HIDDEN slice.  Each thread owns BF16_VEC
     // contiguous positions per step; stride BLOCK_SIZE*BF16_VEC.
     for (int h = h_start + tid * BF16_VEC; h < h_end; h += BLOCK_SIZE * BF16_VEC)
     {
-        // Load x[h..h+BF16_VEC).  Vector width depends on BF16_VEC.
-        float xf[BF16_VEC];
-        if constexpr (BF16_VEC == 8)
-        {
-            uint4 x_raw = *reinterpret_cast<uint4 const*>(&x_in[tok_x + h]);
-            __nv_bfloat162 const* xp = reinterpret_cast<__nv_bfloat162 const*>(&x_raw);
+        // Reduce split-major O_b partials before pmap.
+        float xf[BF16_VEC] = {};
 #pragma unroll
-            for (int v = 0; v < 4; v++)
-            {
-                float2 f = __bfloat1622float2(xp[v]);
-                xf[2 * v + 0] = f.x;
-                xf[2 * v + 1] = f.y;
-            }
-        }
-        else if constexpr (BF16_VEC == 4)
+        for (int xs = 0; xs < X_NUM_SPLITS; ++xs)
         {
-            uint2 x_raw = *reinterpret_cast<uint2 const*>(&x_in[tok_x + h]);
-            __nv_bfloat162 const* xp = reinterpret_cast<__nv_bfloat162 const*>(&x_raw);
+            long long const tok_x
+                = (static_cast<long long>(xs) * gridDim.x + static_cast<long long>(token)) * hidden_size;
+            if constexpr (BF16_VEC == 8)
+            {
+                uint4 x_raw = *reinterpret_cast<uint4 const*>(&x_in[tok_x + h]);
+                __nv_bfloat162 const* xp = reinterpret_cast<__nv_bfloat162 const*>(&x_raw);
 #pragma unroll
-            for (int v = 0; v < 2; v++)
+                for (int v = 0; v < 4; v++)
+                {
+                    float2 f = __bfloat1622float2(xp[v]);
+                    xf[2 * v + 0] += f.x;
+                    xf[2 * v + 1] += f.y;
+                }
+            }
+            else if constexpr (BF16_VEC == 4)
             {
-                float2 f = __bfloat1622float2(xp[v]);
-                xf[2 * v + 0] = f.x;
-                xf[2 * v + 1] = f.y;
+                uint2 x_raw = *reinterpret_cast<uint2 const*>(&x_in[tok_x + h]);
+                __nv_bfloat162 const* xp = reinterpret_cast<__nv_bfloat162 const*>(&x_raw);
+#pragma unroll
+                for (int v = 0; v < 2; v++)
+                {
+                    float2 f = __bfloat1622float2(xp[v]);
+                    xf[2 * v + 0] += f.x;
+                    xf[2 * v + 1] += f.y;
+                }
+            }
+            else
+            { // BF16_VEC == 2
+                unsigned x_raw = *reinterpret_cast<unsigned const*>(&x_in[tok_x + h]);
+                __nv_bfloat162 xp = *reinterpret_cast<__nv_bfloat162 const*>(&x_raw);
+                float2 f = __bfloat1622float2(xp);
+                xf[0] += f.x;
+                xf[1] += f.y;
             }
-        }
-        else
-        { // BF16_VEC == 2
-            unsigned x_raw = *reinterpret_cast<unsigned const*>(&x_in[tok_x + h]);
-            __nv_bfloat162 xp = *reinterpret_cast<__nv_bfloat162 const*>(&x_raw);
-            float2 f = __bfloat1622float2(xp);
-            xf[0] = f.x;
-            xf[1] = f.y;
         }
 
         // Seed new_r[j, v] = pm[j] * x[v] for all j
diff --git a/cpp/tensorrt_llm/thop/mhcOp.cpp b/cpp/tensorrt_llm/thop/mhcOp.cpp
index 484778e05f52..b1d2ef1229a3 100644
--- a/cpp/tensorrt_llm/thop/mhcOp.cpp
+++ b/cpp/tensorrt_llm/thop/mhcOp.cpp
@@ -81,15 +81,25 @@ void mhcFusedHcOp(torch::Tensor x_prev, torch::Tensor residual_prev, torch::Tens
     torch::Tensor y_acc_workspace, torch::Tensor r_acc_workspace, torch::Tensor done_counter_workspace, int64_t M,
     int64_t hidden_size, int64_t hc_mult, double rms_eps, double hc_pre_eps, double hc_sinkhorn_eps,
     double hc_post_mult_value, int64_t sinkhorn_repeat, int64_t backend, int64_t tile_n, int64_t num_k_splits,
-    int64_t bigfuse_block_size, int64_t tile_m, c10::optional<torch::Tensor> norm_weight, double norm_eps)
+    int64_t bigfuse_block_size, int64_t tile_m, c10::optional<torch::Tensor> norm_weight, double norm_eps,
+    int64_t x_num_splits)
 {
     auto stream = at::cuda::getCurrentCUDAStream();
 
-    // Fused next-layer RMSNorm on layer_input_cur. All four backends (Path B/D
-    // for MMA and Path E/F for FMA) support it now: Path B/E fuse the norm into
-    // the bigfuse kernel's Phase 2 layer_input write; Path D/F fuse it into the
-    // single-kernel Phase 4 epilogue. When norm_weight is null, layer_input is
-    // left un-normalized (caller must run RMSNorm separately).
+    TORCH_CHECK(x_prev.dtype() == torch::kBFloat16, "mhc_fused_hc: x_prev must be bfloat16");
+    TORCH_CHECK(x_prev.is_contiguous(), "mhc_fused_hc: x_prev must be contiguous");
+    TORCH_CHECK(x_num_splits == 1 || x_num_splits == 2 || x_num_splits == 4,
+        "mhc_fused_hc: x_num_splits must be 1, 2, or 4, got ", x_num_splits);
+    TORCH_CHECK(x_prev.dim() == 2 && x_prev.size(0) == x_num_splits * M && x_prev.size(1) == hidden_size,
+        "mhc_fused_hc: x_prev must have shape [x_num_splits * M, hidden_size]");
+
+    auto const* x_prev_ptr = reinterpret_cast<__nv_bfloat16 const*>(x_prev.data_ptr<at::BFloat16>());
+    // Half-MMA uses an x-ring; half-FMA reduces in registers.
+    bool const fuse_x_reduction = x_num_splits > 1;
+    TORCH_CHECK(!fuse_x_reduction || backend == 0 || backend == 1,
+        "mhc_fused_hc: split x is supported only by fused_half_mma/fma");
+
+    // A null norm_weight leaves layer_input_cur unnormalized.
     __nv_bfloat16 const* norm_weight_ptr = nullptr;
     if (norm_weight.has_value() && norm_weight->defined())
     {
@@ -107,7 +117,7 @@ void mhcFusedHcOp(torch::Tensor x_prev, torch::Tensor residual_prev, torch::Tens
     //   3 = fused_all_fma  (1-kernel: fused_pmap_gemm_fma_allinone,            Path F)
     if (backend == 3)
     {
-        tk::mhcFusedHcFmaAllInOneLaunch(reinterpret_cast<__nv_bfloat16 const*>(x_prev.data_ptr<at::BFloat16>()),
+        tk::mhcFusedHcFmaAllInOneLaunch(x_prev_ptr,
             reinterpret_cast<__nv_bfloat16 const*>(residual_prev.data_ptr<at::BFloat16>()),
             post_mix_prev.data_ptr<float>(), comb_mix_prev.data_ptr<float>(), w.data_ptr<float>(),
             hc_scale.data_ptr<float>(), hc_base.data_ptr<float>(),
@@ -123,7 +133,7 @@ void mhcFusedHcOp(torch::Tensor x_prev, torch::Tensor residual_prev, torch::Tens
     }
     if (backend == 2)
     {
-        tk::mhcFusedHcAllInOneLaunch(reinterpret_cast<__nv_bfloat16 const*>(x_prev.data_ptr<at::BFloat16>()),
+        tk::mhcFusedHcAllInOneLaunch(x_prev_ptr,
             reinterpret_cast<__nv_bfloat16 const*>(residual_prev.data_ptr<at::BFloat16>()),
             post_mix_prev.data_ptr<float>(), comb_mix_prev.data_ptr<float>(), w.data_ptr<float>(),
             hc_scale.data_ptr<float>(), hc_base.data_ptr<float>(),
@@ -138,7 +148,7 @@ void mhcFusedHcOp(torch::Tensor x_prev, torch::Tensor residual_prev, torch::Tens
     }
     if (backend == 1)
     {
-        tk::mhcFusedHcFmaLaunch(reinterpret_cast<__nv_bfloat16 const*>(x_prev.data_ptr<at::BFloat16>()),
+        tk::mhcFusedHcFmaLaunch(x_prev_ptr,
             reinterpret_cast<__nv_bfloat16 const*>(residual_prev.data_ptr<at::BFloat16>()),
             post_mix_prev.data_ptr<float>(), comb_mix_prev.data_ptr<float>(), w.data_ptr<float>(),
             hc_scale.data_ptr<float>(), hc_base.data_ptr<float>(),
@@ -148,20 +158,21 @@ void mhcFusedHcOp(torch::Tensor x_prev, torch::Tensor residual_prev, torch::Tens
             static_cast<int>(hidden_size), static_cast<int>(hc_mult), static_cast<int>(tile_n),
             static_cast<int>(num_k_splits), static_cast<int>(bigfuse_block_size), static_cast<float>(rms_eps),
             static_cast<float>(hc_pre_eps), static_cast<float>(hc_sinkhorn_eps), static_cast<float>(hc_post_mult_value),
-            static_cast<int>(sinkhorn_repeat), norm_weight_ptr, static_cast<float>(norm_eps), stream);
+            static_cast<int>(sinkhorn_repeat), norm_weight_ptr, static_cast<float>(norm_eps), stream,
+            fuse_x_reduction ? static_cast<int>(x_num_splits) : 1);
         return;
     }
 
-    tk::mhcFusedHcLaunch(reinterpret_cast<__nv_bfloat16 const*>(x_prev.data_ptr<at::BFloat16>()),
-        reinterpret_cast<__nv_bfloat16 const*>(residual_prev.data_ptr<at::BFloat16>()), post_mix_prev.data_ptr<float>(),
-        comb_mix_prev.data_ptr<float>(), w.data_ptr<float>(), hc_scale.data_ptr<float>(), hc_base.data_ptr<float>(),
+    tk::mhcFusedHcLaunch(x_prev_ptr, reinterpret_cast<__nv_bfloat16 const*>(residual_prev.data_ptr<at::BFloat16>()),
+        post_mix_prev.data_ptr<float>(), comb_mix_prev.data_ptr<float>(), w.data_ptr<float>(),
+        hc_scale.data_ptr<float>(), hc_base.data_ptr<float>(),
         reinterpret_cast<__nv_bfloat16*>(residual_cur.data_ptr<at::BFloat16>()), post_mix_cur.data_ptr<float>(),
         comb_mix_cur.data_ptr<float>(), reinterpret_cast<__nv_bfloat16*>(layer_input_cur.data_ptr<at::BFloat16>()),
         y_acc_workspace.data_ptr<float>(), r_acc_workspace.data_ptr<float>(), static_cast<int>(M),
         static_cast<int>(hidden_size), static_cast<int>(hc_mult), static_cast<int>(num_k_splits),
         static_cast<int>(bigfuse_block_size), static_cast<float>(rms_eps), static_cast<float>(hc_pre_eps),
         static_cast<float>(hc_sinkhorn_eps), static_cast<float>(hc_post_mult_value), static_cast<int>(sinkhorn_repeat),
-        norm_weight_ptr, static_cast<float>(norm_eps), stream);
+        norm_weight_ptr, static_cast<float>(norm_eps), stream, fuse_x_reduction ? static_cast<int>(x_num_splits) : 1);
 }
 
 } // anonymous namespace
@@ -210,7 +221,7 @@ TORCH_LIBRARY_FRAGMENT(trtllm, m)
         "float rms_eps, float hc_pre_eps, float hc_sinkhorn_eps, "
         "float hc_post_mult_value, int sinkhorn_repeat, "
         "int backend=0, int tile_n=0, int num_k_splits=0, int bigfuse_block_size=0, "
-        "int tile_m=1, Tensor? norm_weight=None, float norm_eps=0.0) -> ()");
+        "int tile_m=1, Tensor? norm_weight=None, float norm_eps=0.0, int x_num_splits=1) -> ()");
 }
 
 TORCH_LIBRARY_IMPL(trtllm, CUDA, m)
diff --git a/tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py b/tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py
index 17ae63da2a8d..aa76a8643ed2 100644
--- a/tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py
+++ b/tensorrt_llm/_torch/custom_ops/cute_dsl_custom_ops.py
@@ -733,7 +733,7 @@ def forward(
                     max_active_clusters,
                     stream,
                     swap_ab,
-                    options=f"--opt-level 2 --enable-tvm-ffi"
+                    options="--opt-level 2 --enable-tvm-ffi"
                     if self.use_tvm_ffi else "--opt-level 2",
                 )
 
@@ -1141,7 +1141,7 @@ def forward(
                     max_active_clusters,
                     stream,
                     False,  # swap_ab=False for SwiGLU
-                    options=f"--opt-level 2 --enable-tvm-ffi"
+                    options="--opt-level 2 --enable-tvm-ffi"
                     if self.use_tvm_ffi else "--opt-level 2",
                 )
 
@@ -1479,7 +1479,7 @@ def forward(
                     norm_const_cute_tensor,
                     max_active_clusters,
                     stream,
-                    options=f"--opt-level 2 --enable-tvm-ffi"
+                    options="--opt-level 2 --enable-tvm-ffi"
                     if self.use_tvm_ffi else "--opt-level 2",
                 )
 
@@ -3439,7 +3439,7 @@ def forward(
                     c_cute_tensor,
                     max_active_clusters=max_active_clusters,
                     stream=stream,
-                    options=f"--opt-level 2 --enable-tvm-ffi"
+                    options="--opt-level 2 --enable-tvm-ffi"
                     if self.use_tvm_ffi else "--opt-level 2",
                 )
                 self.__class__.kernel_cache[cache_key] = compiled_gemm
@@ -3534,6 +3534,297 @@ def _(
         ret = mat_a.new_empty(shape, dtype=torch.bfloat16)
         return ret
 
+    class CuteDSLFp8SplitKGemmRunner:
+        """DSV4-Pro O_b runner using native block-scaled tcgen05 MMA."""
+
+        kernel_class = Sm100BlockScaledPersistentDenseGemmKernel
+        kernel_cache = dict()
+        alpha_cache = dict()
+        active_clusters_cache = dict()
+        _TILE_M = 256
+        _CLUSTER_M = 2
+        _M_BUCKET = 64
+        _TILE_N_GRANULARITY = 16
+        _MAX_SWAPPED_TILE_N = 240
+        _SCALE_BLOCK_K = 128
+        _SCALES_PER_WORD = 4
+        _PACKED_SCALE_K = _SCALE_BLOCK_K * _SCALES_PER_WORD
+
+        def __init__(self, use_tvm_ffi: bool = True):
+            self.use_tvm_ffi = use_tvm_ffi
+
+        @classmethod
+        def _get_swapped_tile_n(cls, m: int, n: int,
+                                max_active_clusters: int) -> int:
+            bucket_m = pad_up(m, cls._M_BUCKET)
+            if bucket_m <= cls._TILE_M * cls._CLUSTER_M:
+                return pad_up(ceil_div(bucket_m, cls._CLUSTER_M),
+                              cls._TILE_N_GRANULARITY)
+
+            # Fill complete cluster waves within the tile-N cap.
+            clusters_per_n_tile = ceil_div(n, cls._TILE_M)
+            min_n_tiles = ceil_div(bucket_m, cls._MAX_SWAPPED_TILE_N)
+            waves = ceil_div(
+                min_n_tiles * clusters_per_n_tile,
+                max_active_clusters,
+            )
+            available_clusters = waves * max_active_clusters
+            n_tiles = max(
+                1,
+                available_clusters // clusters_per_n_tile,
+            )
+            return pad_up(ceil_div(bucket_m, n_tiles), cls._TILE_N_GRANULARITY)
+
+        @classmethod
+        def _get_no_swap_tile_n(cls, m: int, num_splits: int) -> Optional[int]:
+            if num_splits != 1:
+                return None
+            # Small-M no-swap cases avoid underfilled swapped tiles.
+            if (448 <= m <= cls._TILE_M * cls._CLUSTER_M
+                    and m % cls._M_BUCKET == 0):
+                return 208
+
+            # No-swap resonances repeat every seven M tiles from M=2304.
+            m_tiles = ceil_div(m, cls._TILE_M)
+            first_resonant_m_tiles = 9
+            resonance_period_m_tiles = 7
+            is_resonant = (m_tiles >= first_resonant_m_tiles
+                           and (m_tiles - first_resonant_m_tiles) %
+                           resonance_period_m_tiles == 0)
+            if m % cls._TILE_M == 0 and is_resonant:
+                return 224
+            return None
+
+        @classmethod
+        def _get_tactic(cls, m: int, n: int, num_splits: int,
+                        max_active_clusters: int):
+            if num_splits > 1 and m <= cls._M_BUCKET // 2:
+                # A 128x32 tile exposes more work at tiny M.
+                return ((cls._TILE_M // cls._CLUSTER_M,
+                         cls._M_BUCKET // cls._CLUSTER_M), (cls._CLUSTER_M, 1),
+                        True, None, True)
+            no_swap_tile_n = cls._get_no_swap_tile_n(m, num_splits)
+            if no_swap_tile_n is not None:
+                # Grouped L2 traversal regresses at the M=2304 resonance.
+                return ((cls._TILE_M, no_swap_tile_n), (cls._CLUSTER_M, 1),
+                        False, None, m != 2304)
+
+            tile_n = cls._get_swapped_tile_n(m, n, max_active_clusters)
+            if num_splits > 1 and m <= 2 * cls._M_BUCKET:
+                tile_n = cls._M_BUCKET if m <= cls._M_BUCKET else 2 * cls._M_BUCKET
+            max_ab_stages = 8 if m <= cls._M_BUCKET else None
+            return ((cls._TILE_M, tile_n), (cls._CLUSTER_M, 1), True,
+                    max_ab_stages, True)
+
+        def forward(self, inputs: List[torch.Tensor], num_splits: int) -> None:
+            a_tensor, a_sf_tensor, b_tensor, b_sf_tensor, partials = inputs
+            m, k = a_tensor.shape
+            n = b_tensor.shape[0]
+            c_tmp = partials.permute(1, 2, 0)
+
+            device_key = a_tensor.device
+            max_active_clusters = self.__class__.active_clusters_cache.get(
+                device_key)
+            if max_active_clusters is None:
+                max_active_clusters = cutlass.utils.HardwareInfo(
+                ).get_max_active_clusters(self._CLUSTER_M)
+                self.__class__.active_clusters_cache[
+                    device_key] = max_active_clusters
+            (mma_tiler_mn, cluster_shape_mn, swap_ab, max_ab_stages,
+             l2_swizzle) = self._get_tactic(m, n, num_splits,
+                                            max_active_clusters)
+            # Large M amortizes a dedicated scale-TMA warp.
+            tma_packed_scales = m >= 512
+            if swap_ab:
+                # Swap M/N to expose more output tiles.
+                kernel_m, kernel_n = n, m
+                kernel_a, kernel_b = b_tensor, a_tensor
+                kernel_sfa, kernel_sfb = b_sf_tensor, a_sf_tensor
+            else:
+                kernel_m, kernel_n = m, n
+                kernel_a, kernel_b = a_tensor, b_tensor
+                kernel_sfa, kernel_sfb = a_sf_tensor, b_sf_tensor
+            kernel_sfa_stride = kernel_sfa.stride(1)
+            kernel_sfb_stride = kernel_sfb.stride(1)
+            # Tail tiles require a runtime MMA-N descriptor.
+            dynamic_mma_n = kernel_n % mma_tiler_mn[1] != 0
+            cache_key = (kernel_m, kernel_n, k, num_splits, kernel_sfa_stride,
+                         kernel_sfb_stride, mma_tiler_mn, cluster_shape_mn,
+                         swap_ab, dynamic_mma_n, max_ab_stages, l2_swizzle,
+                         tma_packed_scales, self.use_tvm_ffi)
+
+            a_ptr = make_ptr(
+                cutlass.Float8E4M3FN,
+                kernel_a.data_ptr(),
+                cute.AddressSpace.gmem,
+                assumed_align=16,
+            )
+            b_ptr = make_ptr(
+                cutlass.Float8E4M3FN,
+                kernel_b.data_ptr(),
+                cute.AddressSpace.gmem,
+                assumed_align=16,
+            )
+            a_sf_ptr = make_ptr(
+                cutlass.Uint32,
+                kernel_sfa.data_ptr(),
+                cute.AddressSpace.gmem,
+                assumed_align=16,
+            )
+            b_sf_ptr = make_ptr(
+                cutlass.Uint32,
+                kernel_sfb.data_ptr(),
+                cute.AddressSpace.gmem,
+                assumed_align=16,
+            )
+            c_cute_tensor = cute.runtime.from_dlpack(c_tmp).mark_layout_dynamic(
+                leading_dim=1)
+            alpha = self.__class__.alpha_cache.get(device_key)
+            if alpha is None:
+                alpha = torch.ones((1, ),
+                                   device=device_key,
+                                   dtype=torch.float32)
+                self.__class__.alpha_cache[device_key] = alpha
+            alpha_cute_tensor = cute.runtime.from_dlpack(alpha)
+            stream = (cute.runtime.make_fake_stream(
+                use_tvm_ffi_env_stream=True) if self.use_tvm_ffi else
+                      cuda.CUstream(torch.cuda.current_stream().cuda_stream))
+
+            if cache_key not in self.__class__.kernel_cache:
+                gemm = self.__class__.kernel_class(
+                    32,
+                    mma_tiler_mn,
+                    cluster_shape_mn,
+                    False,
+                    packed_k128_scales=True,
+                    tma_packed_scales=tma_packed_scales,
+                    dynamic_mma_n=dynamic_mma_n,
+                    apply_alpha=False,
+                    max_ab_stages=max_ab_stages,
+                    l2_swizzle=l2_swizzle,
+                )
+                compiled_gemm = cute.compile(
+                    gemm.wrapper_dsv4_splitk_packed_ue8m0,
+                    kernel_m,
+                    kernel_n,
+                    k,
+                    num_splits,
+                    kernel_sfa_stride,
+                    kernel_sfb_stride,
+                    a_ptr,
+                    b_ptr,
+                    a_sf_ptr,
+                    b_sf_ptr,
+                    c_cute_tensor,
+                    alpha_cute_tensor,
+                    max_active_clusters,
+                    stream,
+                    swap_ab,
+                    options=("--opt-level 2 --enable-tvm-ffi"
+                             if self.use_tvm_ffi else "--opt-level 2"),
+                )
+                self.__class__.kernel_cache[cache_key] = compiled_gemm
+            else:
+                compiled_gemm = self.__class__.kernel_cache[cache_key]
+
+            if self.use_tvm_ffi:
+                compiled_gemm(
+                    kernel_a.data_ptr(),
+                    kernel_b.data_ptr(),
+                    kernel_sfa.data_ptr(),
+                    kernel_sfb.data_ptr(),
+                    c_tmp,
+                    alpha,
+                )
+            else:
+                compiled_gemm(
+                    a_ptr,
+                    b_ptr,
+                    a_sf_ptr,
+                    b_sf_ptr,
+                    c_cute_tensor,
+                    alpha_cute_tensor,
+                    stream=stream,
+                )
+
+    @torch.library.custom_op(
+        "trtllm::dsv4_fp8_splitk_gemm",
+        mutates_args=("partials", ),
+        device_types="cuda",
+    )
+    def dsv4_fp8_splitk_gemm(
+        a: torch.Tensor,
+        sfa: torch.Tensor,
+        b: torch.Tensor,
+        sfb: torch.Tensor,
+        partials: torch.Tensor,
+        num_splits: int,
+    ) -> None:
+        """Emit split-major BF16 partials for the DSV4 O_b FP8 GEMM."""
+        if not is_sm_100f():
+            raise ValueError(
+                f"dsv4_fp8_splitk_gemm requires the SM100 family, got {get_sm_version()}"
+            )
+        if num_splits not in (1, 2, 4):
+            raise ValueError(f"num_splits must be 1, 2, or 4, got {num_splits}")
+        if not all(t.is_cuda for t in (a, sfa, b, sfb, partials)):
+            raise ValueError(
+                "all dsv4_fp8_splitk_gemm tensors must be CUDA tensors")
+        if len({t.device for t in (a, sfa, b, sfb, partials)}) != 1:
+            raise ValueError(
+                "all dsv4_fp8_splitk_gemm tensors must be on the same device")
+        if a.dtype != torch.float8_e4m3fn or b.dtype != torch.float8_e4m3fn:
+            raise TypeError("A and B must have dtype torch.float8_e4m3fn")
+        if sfa.dtype != torch.int32 or sfb.dtype != torch.int32:
+            raise TypeError(
+                "SFA and SFB must contain packed int32 UE8M0 scales")
+        if partials.dtype != torch.bfloat16:
+            raise TypeError("partials must have dtype torch.bfloat16")
+        if a.dim() != 2 or b.dim(
+        ) != 2 or not a.is_contiguous() or not b.is_contiguous():
+            raise ValueError("A and B must be contiguous rank-2 tensors")
+
+        m, k = a.shape
+        n = b.shape[0]
+        if b.shape[1] != k:
+            raise ValueError("A and B K dimensions must match")
+        packed_scale_k = CuteDSLFp8SplitKGemmRunner._PACKED_SCALE_K
+        scale_block_k = CuteDSLFp8SplitKGemmRunner._SCALE_BLOCK_K
+        scales_per_word = CuteDSLFp8SplitKGemmRunner._SCALES_PER_WORD
+        if k % (packed_scale_k * num_splits) != 0:
+            raise ValueError(
+                f"K={k} must be divisible by {packed_scale_k}*num_splits")
+        if n % scale_block_k != 0:
+            raise ValueError(f"N={n} must be divisible by {scale_block_k}")
+        if partials.shape != (num_splits, m, n) or not partials.is_contiguous():
+            raise ValueError(
+                f"partials must be contiguous with shape {(num_splits, m, n)}")
+
+        packed_k = k // packed_scale_k
+        if (sfa.dim() != 2 or sfa.shape != (m, packed_k) or sfa.stride(0) != 1
+                or sfa.stride(1) < m or sfa.stride(1) % scales_per_word != 0):
+            raise ValueError(
+                f"SFA must be packed MN-major [M,K/{packed_scale_k}] with "
+                f"{scales_per_word}-aligned K stride")
+        if (sfb.dim() != 2 or sfb.shape != (n, packed_k) or sfb.stride(0) != 1
+                or sfb.stride(1) != n):
+            raise ValueError(
+                f"SFB must be packed MN-major [N,K/{packed_scale_k}]")
+
+        CuteDSLFp8SplitKGemmRunner(use_tvm_ffi=True).forward(
+            [a, sfa, b, sfb, partials], num_splits)
+
+    @torch.library.register_fake("trtllm::dsv4_fp8_splitk_gemm")
+    def _(
+        a: torch.Tensor,
+        sfa: torch.Tensor,
+        b: torch.Tensor,
+        sfb: torch.Tensor,
+        partials: torch.Tensor,
+        num_splits: int,
+    ) -> None:
+        return None
+
     class CuteDSLFp8BlackwellBmmRunner(TunableRunner):
         kernel_class = Sm100BlockwiseGemmKernel
         kernel_cache = dict()
@@ -3753,7 +4044,7 @@ def forward(
                     c_cute_tensor,
                     max_active_clusters=max_active_clusters,
                     stream=stream,
-                    options=f"--opt-level 2 --enable-tvm-ffi"
+                    options="--opt-level 2 --enable-tvm-ffi"
                     if self.use_tvm_ffi else "--opt-level 2",
                 )
                 self.__class__.kernel_cache[cache_key] = compiled_gemm
@@ -3795,6 +4086,135 @@ def forward(
                     stream=stream,
                 )
 
+        def forward_fp8out(
+            self,
+            inputs: List[torch.Tensor],
+            tactic=None,
+        ) -> None:
+            """Run the fixed DSV4 o_a tactic with FP8 and packed-UE8M0 output."""
+            fixed_tactic = (False, (128, 128), (1, 1))
+            use_2cta_instrs, mma_tiler_mn, cluster_shape_mn = (
+                tactic if isinstance(tactic, tuple) else fixed_tactic)
+            if (use_2cta_instrs, mma_tiler_mn,
+                    cluster_shape_mn) != fixed_tactic:
+                raise ValueError("FP8-output BMM supports only the fixed "
+                                 "(1-CTA, 128x128, 1x1 cluster) tactic")
+            if not self.use_tvm_ffi:
+                raise ValueError("FP8-output BMM requires TVM FFI")
+
+            (a_tensor, b_tensor, a_sf_tensor, b_sf_tensor, c_tensor,
+             sf_out_tensor) = inputs
+            batch_size, m, k = a_tensor.shape
+            n = b_tensor.shape[1]
+            aligned_m = pad_up(m, 4)
+            sf_m = aligned_m
+            sf_k = ceil_div(k, 128)
+            sf_n = ceil_div(n, 128)
+            # SMEM cooperation wins through M=32.
+            use_fp8_smem_epilogue = m <= 32
+            fp8_smem_row_iters = ceil_div(m, 16) if use_fp8_smem_epilogue else 1
+
+            if c_tensor.dtype != torch.float8_e4m3fn or c_tensor.shape != (
+                    batch_size, m, n):
+                raise ValueError(
+                    "FP8-output BMM output must have shape [batch, M, N] "
+                    "and dtype float8_e4m3fn")
+            if sf_out_tensor.dtype != torch.int32 or sf_out_tensor.stride() != (
+                    1, aligned_m):
+                raise ValueError(
+                    "FP8-output BMM scales must be int32 with MN-major "
+                    f"stride (1, {aligned_m})")
+
+            c_tmp = c_tensor.permute(1, 2, 0)
+            cache_key = (
+                "fp8out",
+                use_2cta_instrs,
+                mma_tiler_mn,
+                cluster_shape_mn,
+                self.use_tvm_ffi,
+                fp8_smem_row_iters,
+                use_fp8_smem_epilogue,
+            )
+            if cache_key not in self.__class__.kernel_cache:
+                a_ptr = make_ptr(cutlass.Float8E4M3FN,
+                                 a_tensor.data_ptr(),
+                                 cute.AddressSpace.gmem,
+                                 assumed_align=16)
+                b_ptr = make_ptr(cutlass.Float8E4M3FN,
+                                 b_tensor.data_ptr(),
+                                 cute.AddressSpace.gmem,
+                                 assumed_align=16)
+                a_sf_ptr = make_ptr(cutlass.Float32,
+                                    a_sf_tensor.data_ptr(),
+                                    cute.AddressSpace.gmem,
+                                    assumed_align=16)
+                b_sf_ptr = make_ptr(cutlass.Float32,
+                                    b_sf_tensor.data_ptr(),
+                                    cute.AddressSpace.gmem,
+                                    assumed_align=16)
+                c_cute_tensor = cute.runtime.from_dlpack(
+                    c_tmp).mark_layout_dynamic(leading_dim=1)
+                sf_out_ptr = make_ptr(cutlass.Int32,
+                                      sf_out_tensor.data_ptr(),
+                                      cute.AddressSpace.gmem,
+                                      assumed_align=4)
+
+                gemm = self.__class__.kernel_class(
+                    cutlass.Float32,
+                    use_2cta_instrs=use_2cta_instrs,
+                    mma_tiler_mn=mma_tiler_mn,
+                    cluster_shape_mn=cluster_shape_mn,
+                )
+                cluster_size = cluster_shape_mn[0] * cluster_shape_mn[1]
+                max_active_clusters = cutlass.utils.HardwareInfo(
+                ).get_max_active_clusters(cluster_size)
+                stream = cute.runtime.make_fake_stream(
+                    use_tvm_ffi_env_stream=True)
+                compiled_gemm = cute.compile(
+                    gemm.wrapper_fp8sf,
+                    m,
+                    n,
+                    k,
+                    sf_m,
+                    sf_n,
+                    sf_k,
+                    batch_size,
+                    a_ptr,
+                    b_ptr,
+                    a_sf_ptr,
+                    b_sf_ptr,
+                    c_cute_tensor,
+                    max_active_clusters,
+                    stream,
+                    sf_out_ptr,
+                    aligned_m,
+                    sf_n,
+                    fp8_smem_row_iters,
+                    use_fp8_smem_epilogue,
+                    options="--opt-level 2 --enable-tvm-ffi",
+                )
+                self.__class__.kernel_cache[cache_key] = compiled_gemm
+            else:
+                compiled_gemm = self.__class__.kernel_cache[cache_key]
+
+            compiled_gemm(
+                m,
+                n,
+                k,
+                sf_m,
+                sf_n,
+                sf_k,
+                batch_size,
+                a_tensor.data_ptr(),
+                b_tensor.data_ptr(),
+                a_sf_tensor.data_ptr(),
+                b_sf_tensor.data_ptr(),
+                c_tmp,
+                sf_out_tensor.data_ptr(),
+                aligned_m,
+                sf_n,
+            )
+
     # a/b: fp8, scale: fp32, output: bf16
     @torch.library.custom_op("trtllm::cute_dsl_fp8_bmm_blackwell",
                              mutates_args=("output", ),
@@ -3849,6 +4269,42 @@ def _(
         assert output.shape == (batch_size, m,
                                 n), "CuTe DSL fp8 bmm output shape is incorrect"
 
+    # Emit FP8 O_a output and packed UE8M0 scales in one pass.
+    @torch.library.custom_op("trtllm::cute_dsl_fp8_bmm_blackwell_fp8out",
+                             mutates_args=("output_fp8", "sf_out"),
+                             device_types="cuda")
+    def cute_dsl_fp8_bmm_blackwell_fp8out(
+        input: torch.Tensor,
+        weight: torch.Tensor,
+        input_scale: torch.Tensor,
+        weight_scale: torch.Tensor,
+        output_fp8: torch.Tensor,
+        sf_out: torch.Tensor,
+    ) -> None:
+        """Write [L,M,N] FP8 output and MN-major packed scales."""
+        if not is_sm_100f():
+            raise ValueError(
+                f"cute_dsl_fp8_bmm_blackwell_fp8out requires SM100 family, got {get_sm_version()}"
+            )
+        runner = CuteDSLFp8BlackwellBmmRunner(use_tvm_ffi=True)
+        runner.forward_fp8out(
+            [input, weight, input_scale, weight_scale, output_fp8, sf_out])
+
+    @torch.library.register_fake("trtllm::cute_dsl_fp8_bmm_blackwell_fp8out")
+    def _(
+        input: torch.Tensor,
+        weight: torch.Tensor,
+        input_scale: torch.Tensor,
+        weight_scale: torch.Tensor,
+        output_fp8: torch.Tensor,
+        sf_out: torch.Tensor,
+    ) -> None:
+        L, M, K = input.shape
+        N = weight.shape[1]
+        assert output_fp8.dtype == torch.float8_e4m3fn and output_fp8.shape == (
+            L, M, N)
+        assert sf_out.dtype == torch.int32
+
     # =============================================================================
     # Dense GEMM with SwiGLU Fusion (FC1 Kernel for MoE as Dense GEMM)
     # =============================================================================
diff --git a/tensorrt_llm/_torch/cute_dsl_kernels/blackwell/blockwise_gemm/blockwise_gemm.py b/tensorrt_llm/_torch/cute_dsl_kernels/blackwell/blockwise_gemm/blockwise_gemm.py
index 6881dbd3905f..2fd028f1f793 100644
--- a/tensorrt_llm/_torch/cute_dsl_kernels/blackwell/blockwise_gemm/blockwise_gemm.py
+++ b/tensorrt_llm/_torch/cute_dsl_kernels/blackwell/blockwise_gemm/blockwise_gemm.py
@@ -44,7 +44,7 @@
 # This file is copied and modified from cutlass example https://github.com/NVIDIA/cutlass/blob/main/examples/python/CuTeDSL/blackwell/blockwise_gemm/blockwise_gemm.py
 
 import math
-from typing import Tuple, Type, Union
+from typing import Optional, Tuple, Type, Union
 
 import cuda.bindings.driver as cuda
 import cutlass
@@ -52,9 +52,136 @@
 import cutlass.pipeline as pipeline
 import cutlass.utils as utils
 import cutlass.utils.blackwell_helpers as sm100_utils
+from cutlass._mlir.dialects import llvm
 from cutlass.cute.nvgpu import cpasync, tcgen05
+from cutlass.cutlass_dsl import dsl_user_op
 from cutlass.pipeline import pipeline_init_arrive, pipeline_init_wait
 
+_FP8_E4M3_MAX = 448.0
+_MIN_AMAX = 1e-4
+
+
+@dsl_user_op
+def ceil_to_ue8m0_device(
+    amax_f32: cutlass.Float32,
+    *,
+    loc=None,
+    ip=None,
+) -> cutlass.Float32:
+    """Round an FP32 scale up to UE8M0."""
+    i32_ty = cutlass.Int32.mlir_type
+    f32_ty = cutlass.Float32.mlir_type
+    bits = cutlass.Int32(
+        llvm.bitcast(i32_ty, cutlass.Float32(amax_f32).ir_value(loc=loc, ip=ip), loc=loc, ip=ip)
+    )
+    exp = (bits >> cutlass.Int32(23)) & cutlass.Int32(0xFF)
+    # A nonzero mantissa rounds the exponent up.
+    mantissa = bits & cutlass.Int32(0x7FFFFF)
+    mantissa_nonzero = (mantissa + cutlass.Int32(0x7FFFFF)) >> cutlass.Int32(23)
+    exp = exp + mantissa_nonzero
+    # Clamp away zero and reserved exponents.
+    exp = cutlass.Int32(
+        llvm.inline_asm(
+            i32_ty,
+            [exp.ir_value(loc=loc, ip=ip)],
+            "max.s32 $0, $1, 1;",
+            "=r,r",
+            has_side_effects=False,
+            loc=loc,
+            ip=ip,
+        )
+    )
+    exp = cutlass.Int32(
+        llvm.inline_asm(
+            i32_ty,
+            [exp.ir_value(loc=loc, ip=ip)],
+            "min.s32 $0, $1, 254;",
+            "=r,r",
+            has_side_effects=False,
+            loc=loc,
+            ip=ip,
+        )
+    )
+    sf_bits = exp << cutlass.Int32(23)
+    return cutlass.Float32(llvm.bitcast(f32_ty, sf_bits.ir_value(loc=loc, ip=ip), loc=loc, ip=ip))
+
+
+@dsl_user_op
+def stg_u8_raw(
+    addr: cutlass.Int64,
+    value: cutlass.Int32,
+    *,
+    loc=None,
+    ip=None,
+) -> None:
+    """Store the low byte of value to a global address."""
+    from cutlass._mlir.dialects import llvm as _llvm
+
+    _llvm.inline_asm(
+        None,
+        [
+            cutlass.Int64(addr).ir_value(loc=loc, ip=ip),
+            cutlass.Int32(value).ir_value(loc=loc, ip=ip),
+        ],
+        "st.global.u8 [$0], $1;",
+        "l,r",
+        has_side_effects=True,
+        loc=loc,
+        ip=ip,
+    )
+
+
+@dsl_user_op
+def abs_f32_device(
+    value: cutlass.Float32,
+    *,
+    loc=None,
+    ip=None,
+) -> cutlass.Float32:
+    return cutlass.Float32(
+        llvm.inline_asm(
+            cutlass.Float32.mlir_type,
+            [cutlass.Float32(value).ir_value(loc=loc, ip=ip)],
+            "abs.f32 $0, $1;",
+            "=f,f",
+            has_side_effects=False,
+            loc=loc,
+            ip=ip,
+        )
+    )
+
+
+@dsl_user_op
+def max_nan_f32_device(
+    lhs: cutlass.Float32,
+    rhs: cutlass.Float32,
+    *,
+    loc=None,
+    ip=None,
+) -> cutlass.Float32:
+    return cutlass.Float32(
+        llvm.inline_asm(
+            cutlass.Float32.mlir_type,
+            [
+                cutlass.Float32(lhs).ir_value(loc=loc, ip=ip),
+                cutlass.Float32(rhs).ir_value(loc=loc, ip=ip),
+            ],
+            "max.NaN.f32 $0, $1, $2;",
+            "=f,f,f",
+            has_side_effects=False,
+            loc=loc,
+            ip=ip,
+        )
+    )
+
+
+def _tree_reduce_max_nan_f32(values):
+    """Build a balanced max tree for a trace-time value list."""
+    while len(values) > 1:
+        values = [max_nan_f32_device(values[i], values[i + 1]) for i in range(0, len(values), 2)]
+    return values[0]
+
+
 """
 High-performance persistent blockwise dense GEMM (C = (SFA * A) * (SFB * B)) example for the NVIDIA Blackwell
 architecture using CUTE DSL.
@@ -165,6 +292,7 @@ def __init__(
         use_2cta_instrs: bool,
         mma_tiler_mn: Tuple[int, int],
         cluster_shape_mn: Tuple[int, int],
+        smem_reserve_bytes: int = 1024,
     ):
         """Initializes the configuration for a Blackwell blockwise dense GEMM kernel.
 
@@ -187,11 +315,15 @@ def __init__(
         :type use_2cta_instrs: bool
         :param cluster_shape_mn: Tuple (ClusterM, ClusterN) shape of the cluster.
         :type cluster_shape_mn: Tuple[int, int]
+        :param smem_reserve_bytes: Shared-memory bytes reserved for barriers,
+            alignment, and other non-tensor storage.
+        :type smem_reserve_bytes: int
         """
 
         self.acc_dtype: Type[cutlass.Numeric] = acc_dtype
         self.use_2cta_instrs = use_2cta_instrs
         self.cluster_shape_mn = cluster_shape_mn
+        self.smem_reserve_bytes = smem_reserve_bytes
         # K dimension is deferred in _setup_attributes
         self.mma_tiler = (*mma_tiler_mn, 1)
 
@@ -341,7 +473,11 @@ def _setup_attributes(self):
             self.scale_n_per_tile * self.scale_k_per_tile,
             self.num_smem_capacity,
             self.occupancy,
+            self.smem_reserve_bytes,
         )
+        if self.fp8_smem_epi_mode:
+            # Trade one A/B stage for the BF16 epilogue tile.
+            self.num_ab_stage -= 1
 
         # Compute A/B/C/Scale shared memory layout
         self.a_smem_layout_staged = sm100_utils.make_smem_layout_a(
@@ -362,6 +498,13 @@ def _setup_attributes(self):
             self.epi_tile,
             self.num_c_stage,
         )
+        self.fp8_bf16_stage_count = self.cta_tile_shape_mnk[1] // cute.size(self.epi_tile[1].shape)
+        self.fp8_bf16_smem_layout_staged = sm100_utils.make_smem_layout_epi(
+            cutlass.BFloat16,
+            self.c_layout,
+            self.epi_tile,
+            self.fp8_bf16_stage_count,
+        )
         self.sfa_smem_layout_staged = cute.make_layout(
             (
                 (self.scale_granularity_m, self.scale_m_per_tile),
@@ -401,6 +544,11 @@ def __call__(
         max_active_clusters: cutlass.Constexpr,
         stream: cuda.CUstream,
         epilogue_op: cutlass.Constexpr = lambda x: x,
+        sf_out_tensor: Optional[cute.Tensor] = None,
+        sf_aligned_mn: int = 0,
+        n_tiles_per_group: int = 0,
+        fp8_smem_row_iters: cutlass.Constexpr = 1,
+        use_fp8_smem_epilogue: cutlass.Constexpr = False,
     ):
         """Execute the GEMM operation in steps:
         - Setup static attributes before smem/grid/tma computation
@@ -425,6 +573,18 @@ def __call__(
         :type stream: cuda.CUstream
         :param epilogue_op: Optional elementwise lambda function to apply to the output tensor
         :type epilogue_op: cutlass.Constexpr
+        :param sf_out_tensor: cute.Tensor pointing to the scale-factor output buffer, or None for BF16 mode.
+        :type sf_out_tensor: Optional[cute.Tensor]
+        :param sf_aligned_mn: MN-major stride (in int32 elements) of the scale-factor buffer, or 0 for BF16 mode.
+        :type sf_aligned_mn: int
+        :param n_tiles_per_group: Number of N-tiles per group (per-group ceil_div(N, BLOCK_N)). Used in batched FP8+SF
+            mode to compute the correct flattened kblock index: global_kblock = group * n_tiles_per_group + tile_n.
+            Pass 0 (or omit) for BF16 mode or G=1 FP8 mode.
+        :type n_tiles_per_group: int
+        :param fp8_smem_row_iters: Number of 16-row groups processed by the cooperative FP8 epilogue.
+        :type fp8_smem_row_iters: int
+        :param use_fp8_smem_epilogue: Enable the BF16-SMEM cooperative FP8 epilogue.
+        :type use_fp8_smem_epilogue: bool
         :raises TypeError: If input data types are incompatible with the MMA instruction.
         """
         # Setup static attributes before smem/grid/tma computation
@@ -437,6 +597,11 @@ def __call__(
         self.b_major_mode = utils.LayoutEnum.from_tensor(b).mma_major_mode()
         self.c_layout = utils.LayoutEnum.from_tensor(c)
 
+        # Select the FP8 epilogue at trace time.
+        self.fp8_sf_mode = sf_out_tensor is not None
+        self.fp8_smem_epi_mode = self.fp8_sf_mode and use_fp8_smem_epilogue
+        self.fp8_smem_row_iters = fp8_smem_row_iters
+
         # Check if input data types are compatible with MMA instruction
         if cutlass.const_expr(self.a_dtype != self.b_dtype):
             raise TypeError(f"Type must match: {self.a_dtype} != {self.b_dtype}")
@@ -535,6 +700,9 @@ def __call__(
         self.buffer_align_bytes = 1024
 
         c_smem_size = cute.cosize(self.c_smem_layout_staged.outer)
+        fp8_bf16_smem_size = (
+            cute.cosize(self.fp8_bf16_smem_layout_staged.outer) if self.fp8_smem_epi_mode else 0
+        )
 
         # Define shared storage for kernel
         @cute.struct
@@ -560,6 +728,11 @@ class SharedStorage:
                 ],
                 self.buffer_align_bytes,
             ]
+            # BF16 staging for cooperative FP8 quantization.
+            sEpiBf16: cute.struct.Align[
+                cute.struct.MemRange[cutlass.BFloat16, fp8_bf16_smem_size],
+                self.buffer_align_bytes,
+            ]
             # (MMA, MMA_M, MMA_K, STAGE)
             sA: cute.struct.Align[
                 cute.struct.MemRange[self.a_dtype, cute.cosize(self.a_smem_layout_staged.outer)],
@@ -583,6 +756,43 @@ class SharedStorage:
 
         self.shared_storage = SharedStorage
 
+        # FP8 output always requires a scale buffer.
+        if cutlass.const_expr(self.fp8_sf_mode and self.c_dtype is not cutlass.Float8E4M3FN):
+            raise ValueError(
+                f"fp8+SF mode requires c_dtype=Float8E4M3FN, got {self.c_dtype}. "
+                "Pass sf_out_tensor=None for non-fp8 output."
+            )
+        if cutlass.const_expr(not self.fp8_sf_mode and self.c_dtype is cutlass.Float8E4M3FN):
+            raise ValueError(
+                "c_dtype=Float8E4M3FN requires sf_out_tensor to be provided. "
+                "Direct fp8 output without a scale-factor buffer is not supported."
+            )
+
+        # The FP8 lane mapping is fixed to one 128x128 CTA tile.
+        if cutlass.const_expr(self.fp8_sf_mode and self.use_2cta_instrs):
+            raise ValueError(
+                "fp8+SF mode does not support use_2cta_instrs=True. "
+                "Use single-CTA mode (use_2cta_instrs=False)."
+            )
+        if cutlass.const_expr(
+            self.fp8_sf_mode
+            and (self.cta_tile_shape_mnk[0] != 128 or self.cta_tile_shape_mnk[1] != 128)
+        ):
+            raise ValueError(
+                f"fp8+SF mode requires cta_tile_shape_mnk[:2]==(128,128), "
+                f"got {self.cta_tile_shape_mnk[:2]}. "
+                "Use mma_tiler_mn=(128,128) with single-CTA."
+            )
+        if cutlass.const_expr(self.fp8_sf_mode and len(self.epilog_warp_id) != 4):
+            raise ValueError(
+                f"fp8+SF mode requires exactly 4 epilogue warps, got {len(self.epilog_warp_id)}."
+            )
+
+        # Keep the scale pointer as a tensor for CuTe ABI compatibility.
+        kernel_sf_out_tensor = sf_out_tensor
+        kernel_sf_aligned_mn = cutlass.Int32(sf_aligned_mn)
+        kernel_n_tiles_per_group = cutlass.Int32(n_tiles_per_group)
+
         # Launch the kernel synchronously
         self.kernel(
             tiled_mma,
@@ -598,11 +808,15 @@ class SharedStorage:
             self.a_smem_layout_staged,
             self.b_smem_layout_staged,
             self.c_smem_layout_staged,
+            self.fp8_bf16_smem_layout_staged,
             self.sfa_smem_layout_staged,
             self.sfb_smem_layout_staged,
             self.epi_tile,
             self.tile_sched_params,
             epilogue_op,
+            kernel_sf_out_tensor,
+            kernel_sf_aligned_mn,
+            kernel_n_tiles_per_group,
         ).launch(
             grid=grid,
             block=[self.threads_per_cta, 1, 1],
@@ -630,11 +844,15 @@ def kernel(
         a_smem_layout_staged: cute.ComposedLayout,
         b_smem_layout_staged: cute.ComposedLayout,
         c_smem_layout_staged: Union[cute.Layout, cute.ComposedLayout, None],
+        fp8_bf16_smem_layout_staged: cute.ComposedLayout,
         sfa_smem_layout_staged: cute.Layout,
         sfb_smem_layout_staged: cute.Layout,
         epi_tile: cute.Tile,
         tile_sched_params: utils.PersistentTileSchedulerParams,
         epilogue_op: cutlass.Constexpr,
+        mSfOut: Optional[cute.Tensor],
+        sf_aligned_mn: cutlass.Int32,
+        n_tiles_per_group: cutlass.Int32,
     ):
         """
         GPU device kernel performing the Persistent batched GEMM computation.
@@ -642,6 +860,9 @@ def kernel(
         warp_idx = cute.arch.warp_idx()
         warp_idx = cute.arch.make_warp_uniform(warp_idx)
         lane_idx = cute.arch.lane_idx()
+        # Recover the scale address from the tensor iterator.
+        if cutlass.const_expr(self.fp8_sf_mode):
+            sf_out_addr = mSfOut.iterator.toint()
 
         #
         # Prefetch tma desc
@@ -773,6 +994,12 @@ def kernel(
         #
         # (EPI_TILE_M, EPI_TILE_N, STAGE)
         sC = storage.sC.get_tensor(c_smem_layout_staged.outer, swizzle=c_smem_layout_staged.inner)
+        sEpiBf16 = None
+        if cutlass.const_expr(self.fp8_smem_epi_mode):
+            sEpiBf16 = storage.sEpiBf16.get_tensor(
+                fp8_bf16_smem_layout_staged.outer,
+                swizzle=fp8_bf16_smem_layout_staged.inner,
+            )
         # (MMA, MMA_M, MMA_K, STAGE)
         sA = storage.sA.get_tensor(a_smem_layout_staged.outer, swizzle=a_smem_layout_staged.inner)
         # (MMA, MMA_N, MMA_K, STAGE)
@@ -903,8 +1130,13 @@ def kernel(
         # Partition global/shared tensor for TMA load A/B
         #
         # load scaleA/scaleB
+        scale_copy_op = (
+            cute.nvgpu.CopyUniversalOp()
+            if mSFA_mkl.element_type.width < 32
+            else cute.nvgpu.cpasync.CopyG2SOp()
+        )
         atom_copy = cute.make_copy_atom(
-            cute.nvgpu.cpasync.CopyG2SOp(),
+            scale_copy_op,
             mSFA_mkl.element_type,
             num_bits_per_copy=mSFA_mkl.element_type.width,
         )
@@ -1491,11 +1723,11 @@ def kernel(
 
                 tTR_rSFA = cute.make_rmem_tensor(
                     cute.slice_(tTR_sSFA, (None, None, None, 0, None, 0)).shape,
-                    self.acc_dtype,
+                    self.sfa_dtype,
                 )
                 tTR_rSFB = cute.make_rmem_tensor(
                     cute.slice_(tTR_sSFB, (None, None, None, 0, None, 0)).shape,
-                    self.acc_dtype,
+                    self.sfb_dtype,
                 )
 
                 scale_consumer_state.reset_count()
@@ -1529,14 +1761,19 @@ def kernel(
                         (None, None, None, 0, None, scale_consumer_state.index),
                     )
 
-                    scale_atom_copy = cute.make_copy_atom(
+                    sfa_atom_copy = cute.make_copy_atom(
+                        cute.nvgpu.CopyUniversalOp(),
+                        self.sfa_dtype,
+                        num_bits_per_copy=self.sfa_dtype.width,
+                    )
+                    sfb_atom_copy = cute.make_copy_atom(
                         cute.nvgpu.CopyUniversalOp(),
-                        self.acc_dtype,
-                        num_bits_per_copy=self.acc_dtype.width,
+                        self.sfb_dtype,
+                        num_bits_per_copy=self.sfb_dtype.width,
                     )
 
-                    cute.copy(scale_atom_copy, tTR_sSFA_slice, tTR_rSFA)
-                    cute.copy(scale_atom_copy, tTR_sSFB_slice, tTR_rSFB)
+                    cute.copy(sfa_atom_copy, tTR_sSFA_slice, tTR_rSFA)
+                    cute.copy(sfb_atom_copy, tTR_sSFB_slice, tTR_rSFB)
 
                     #
                     # Wait for accumulator buffer full
@@ -1566,8 +1803,8 @@ def kernel(
 
                         acc_vec = tTR_rAcc.load()
                         final_vec = tTR_rAcc_subtile.load()
-                        scale_a = tTR_rSFA_subtile.load()
-                        scale_b = tTR_rSFB_subtile.load()
+                        scale_a = tTR_rSFA_subtile.load().to(self.acc_dtype)
+                        scale_b = tTR_rSFB_subtile.load().to(self.acc_dtype)
                         scale = scale_a * scale_b
                         final_vec = acc_vec * scale + final_vec
                         tTR_rAcc_subtile.store(final_vec.to(self.acc_dtype))
@@ -1671,6 +1908,21 @@ def kernel(
             tiled_copy_r2s, tRS_rC, tRS_sC = self.epilog_smem_copy_and_partition(
                 tiled_copy_t2r, tTR_rC, epi_tidx, sC
             )
+            tiled_copy_r2s_bf16 = None
+            tRS_rBf16 = None
+            tRS_sBf16 = None
+            if cutlass.const_expr(self.fp8_smem_epi_mode):
+                tTR_rBf16 = cute.make_rmem_tensor(tTR_rAcc.shape, cutlass.BFloat16)
+                (
+                    tiled_copy_r2s_bf16,
+                    tRS_rBf16,
+                    tRS_sBf16,
+                ) = self.epilog_bf16_smem_copy_and_partition(
+                    tiled_copy_t2r,
+                    tTR_rBf16,
+                    epi_tidx,
+                    sEpiBf16,
+                )
             (
                 tma_atom_c,
                 bSG_sC,
@@ -1716,6 +1968,7 @@ def kernel(
             is_valid_tile = tile_info[3] == 1
 
             num_prev_subtiles = cutlass.Int32(0)
+            sf_base_offset = cutlass.Int32(0)
 
             while is_valid_tile:
                 mma_tile_coord_mnl = (
@@ -1755,53 +2008,281 @@ def kernel(
                 # Store accumulator to global memory in subtiles
                 #
                 subtile_cnt = cute.size(tTR_tAcc.shape, mode=[3])
-                for subtile_idx in cutlass.range(subtile_cnt):
-                    #
-                    # Load accumulator from tensor memory buffer to register
-                    #
-                    tTR_tAcc_mn = tTR_tAcc[(None, None, None, subtile_idx)]
-                    cute.copy(tiled_copy_t2r, tTR_tAcc_mn, tTR_rAcc)
 
-                    #
-                    # Convert to C type
-                    #
-                    acc_vec = tiled_copy_r2s.retile(tTR_rAcc).load()
-                    acc_vec = epilogue_op(acc_vec.to(self.c_dtype))
-                    tRS_rC.store(acc_vec)
+                if cutlass.const_expr(self.fp8_sf_mode and not self.fp8_smem_epi_mode):
+                    # Register FP8 epilogue: one thread owns one M row.
+                    _num_elems_per_subtile = cute.size(tTR_rAcc.shape)
+                    _total_elems = subtile_cnt * _num_elems_per_subtile
+                    acc_saved_bf16 = cute.make_rmem_tensor((_total_elems,), cutlass.BFloat16)
 
-                    #
-                    # Store C to shared memory
-                    #
-                    num_prev_subtiles = num_prev_subtiles + 1
-                    c_buffer = num_prev_subtiles % self.num_c_stage
-                    cute.copy(
-                        tiled_copy_r2s,
-                        tRS_rC,
-                        tRS_sC[(None, None, None, c_buffer)],
+                    thread_amax = cutlass.Float32(0.0)
+                    for subtile_idx in cutlass.range(subtile_cnt):
+                        tTR_tAcc_mn = tTR_tAcc[(None, None, None, subtile_idx)]
+                        cute.copy(tiled_copy_t2r, tTR_tAcc_mn, tTR_rAcc)
+                        # Retile one TMEM subtile for reduction.
+                        acc_fp32 = tiled_copy_r2s.retile(tTR_rAcc).load()
+                        num_elems = cute.size(acc_fp32.shape)
+                        base = subtile_idx * num_elems
+                        # Trace-time list construction fully unrolls the tree.
+                        _subtile_abs = []
+                        for ei in cutlass.range_constexpr(num_elems):
+                            val_bf16 = cutlass.BFloat16(acc_fp32[ei])
+                            # Cache BF16 to avoid a second TMEM load.
+                            acc_saved_bf16[base + ei] = val_bf16
+                            _subtile_abs.append(abs_f32_device(cutlass.Float32(val_bf16)))
+                        subtile_max = _tree_reduce_max_nan_f32(_subtile_abs)
+                        # Merge the four N subtiles.
+                        thread_amax = max_nan_f32_device(thread_amax, subtile_max)
+
+                    # Avoid zero UE8M0 scales.
+                    thread_amax = cutlass.max(thread_amax, cutlass.Float32(_MIN_AMAX))
+                    sf_fp32 = ceil_to_ue8m0_device(thread_amax / cutlass.Float32(_FP8_E4M3_MAX))
+
+                    # UE8M0 stores only the exponent byte.
+                    i32_ty = cutlass.Int32.mlir_type
+                    sf_bits = cutlass.Int32(llvm.bitcast(i32_ty, sf_fp32.ir_value()))
+                    ue8m0_byte_i32 = (sf_bits >> cutlass.Int32(23)) & cutlass.Int32(0xFF)
+
+                    # Map each epilogue lane to one global M row.
+                    epi_warp_local = warp_idx - cutlass.Int32(self.epilog_warp_id[0])
+                    local_m = epi_warp_local * cutlass.Int32(32) + lane_idx
+                    global_m = (
+                        mma_tile_coord_mnl[0] * cutlass.Int32(self.cta_tile_shape_mnk[0]) + local_m
+                    )
+                    # Flatten group and N tile into the scale K block.
+                    global_kblock = (
+                        mma_tile_coord_mnl[2] * n_tiles_per_group + mma_tile_coord_mnl[1]
                     )
-                    # Fence and barrier to make sure shared memory store is visible to TMA store
+                    # Only the scale store needs an M-tail guard.
+                    runtime_m = cutlass.Int32(mC_mnl.shape[0])
+                    if global_m < runtime_m:
+                        col = global_kblock >> cutlass.Int32(2)
+                        byte_in_col = global_kblock & cutlass.Int32(3)
+                        byte_offset = (
+                            col * sf_aligned_mn * cutlass.Int32(4)
+                            + global_m * cutlass.Int32(4)
+                            + byte_in_col
+                        )
+                        stg_u8_raw(sf_out_addr + cutlass.Int64(byte_offset), ue8m0_byte_i32)
+
+                    # Quantize from the cached BF16 row.
+                    sf_rcp = cutlass.Float32(1.0) / sf_fp32
+                    for subtile_idx in cutlass.range(subtile_cnt):
+                        num_elems = _num_elems_per_subtile
+                        base = subtile_idx * num_elems
+                        acc_scaled = cute.make_rmem_tensor((num_elems,), cutlass.Float32)
+                        for ei in cutlass.range_constexpr(num_elems):
+                            val_bf16 = acc_saved_bf16[base + ei]
+                            acc_scaled[ei] = cutlass.Float32(val_bf16) * sf_rcp
+
+                        tRS_rC.store(acc_scaled.load().to(self.c_dtype))
+
+                        # Stage FP8 output for TMA.
+                        num_prev_subtiles = num_prev_subtiles + 1
+                        c_buffer = num_prev_subtiles % self.num_c_stage
+                        cute.copy(
+                            tiled_copy_r2s,
+                            tRS_rC,
+                            tRS_sC[(None, None, None, c_buffer)],
+                        )
+                        cute.arch.fence_view_async_shared()
+                        self.epilog_sync_barrier.arrive_and_wait()
+
+                        if warp_idx == self.epilog_warp_id[0]:
+                            cute.copy(
+                                tma_atom_c,
+                                bSG_sC[(None, c_buffer)],
+                                bSG_gC[(None, subtile_idx)],
+                            )
+                            c_pipeline.producer_commit()
+                            c_pipeline.producer_acquire()
+                        self.epilog_sync_barrier.arrive_and_wait()
+
+                elif cutlass.const_expr(self.fp8_smem_epi_mode):
+                    # Drain TMEM before cooperative FP8 quantization.
+                    for subtile_idx in cutlass.range(subtile_cnt):
+                        tTR_tAcc_mn = tTR_tAcc[(None, None, None, subtile_idx)]
+                        cute.copy(tiled_copy_t2r, tTR_tAcc_mn, tTR_rAcc)
+                        acc_vec = tiled_copy_r2s_bf16.retile(tTR_rAcc).load()
+                        tRS_rBf16.store(acc_vec.to(cutlass.BFloat16))
+                        cute.copy(
+                            tiled_copy_r2s_bf16,
+                            tRS_rBf16,
+                            tRS_sBf16[(None, None, None, subtile_idx)],
+                        )
+
                     cute.arch.fence_view_async_shared()
                     self.epilog_sync_barrier.arrive_and_wait()
+                    epi_pipeline.consumer_release(epi_consumer_state)
+                    epi_consumer_state.advance()
 
-                    #
-                    # TMA store C to global memory
-                    #
-                    if warp_idx == self.epilog_warp_id[0]:
+                    # Finish prior TMA stores before reusing SMEM stages.
+                    c_pipeline.producer_tail()
+                    self.epilog_sync_barrier.arrive_and_wait()
+
+                    epi_warp_local = warp_idx - cutlass.Int32(self.epilog_warp_id[0])
+                    row_group = lane_idx >> cutlass.Int32(3)
+                    lane_in_group = lane_idx & cutlass.Int32(7)
+                    runtime_m = cutlass.Int32(mC_mnl.shape[0])
+                    global_kblock = (
+                        mma_tile_coord_mnl[2] * n_tiles_per_group + mma_tile_coord_mnl[1]
+                    )
+                    sf_col = global_kblock >> cutlass.Int32(2)
+                    sf_byte_in_col = global_kblock & cutlass.Int32(3)
+                    sf_base_offset = sf_col * sf_aligned_mn * cutlass.Int32(4) + sf_byte_in_col
+                    bf16_copy_atom_s2r = cute.make_copy_atom(
+                        cute.nvgpu.CopyUniversalOp(),
+                        cutlass.BFloat16,
+                        num_bits_per_copy=128,
+                    )
+                    bf16_tiled_copy_s2r = cute.make_tiled_copy_tv(
+                        bf16_copy_atom_s2r,
+                        cute.make_layout((8,)),
+                        cute.make_layout((16,)),
+                    )
+                    bf16_thr_copy_s2r = bf16_tiled_copy_s2r.get_slice(lane_in_group)
+                    fp8_copy_atom_r2s = cute.make_copy_atom(
+                        cute.nvgpu.CopyUniversalOp(),
+                        self.c_dtype,
+                        num_bits_per_copy=128,
+                    )
+                    fp8_tiled_copy_r2s = cute.make_tiled_copy_tv(
+                        fp8_copy_atom_r2s,
+                        cute.make_layout((2,)),
+                        cute.make_layout((16,)),
+                    )
+                    fp8_thr_copy_r2s = fp8_tiled_copy_r2s.get_slice(
+                        lane_in_group & cutlass.Int32(1)
+                    )
+
+                    # Four warps form 16 independent 8-lane row groups.
+                    tile_m_offset = mma_tile_coord_mnl[0] * cutlass.Int32(
+                        self.cta_tile_shape_mnk[0]
+                    )
+                    for row_iter in cutlass.range_constexpr(self.fp8_smem_row_iters):
+                        local_row = (
+                            row_iter * cutlass.Int32(16)
+                            + epi_warp_local * cutlass.Int32(4)
+                            + row_group
+                        )
+                        global_m = tile_m_offset + local_row
+                        row_values = cute.make_rmem_tensor((16,), cutlass.BFloat16)
+                        flat_col_base = lane_in_group * cutlass.Int32(16)
+                        bf16_row = cute.group_modes(
+                            sEpiBf16[(local_row, None, None)],
+                            0,
+                            2,
+                        )
+                        bf16_source = bf16_thr_copy_s2r.partition_S(bf16_row)
+                        bf16_fragment = cute.make_fragment_like(bf16_source)
                         cute.copy(
-                            tma_atom_c,
-                            bSG_sC[(None, c_buffer)],
-                            bSG_gC[(None, subtile_idx)],
+                            bf16_copy_atom_s2r,
+                            bf16_source,
+                            bf16_fragment,
                         )
-                        # Fence and barrier to make sure shared memory store is visible to TMA store
-                        c_pipeline.producer_commit()
-                        c_pipeline.producer_acquire()
+                        row_values.store(bf16_fragment.load())
+
+                        level = []
+                        for value_idx in cutlass.range_constexpr(16):
+                            level.append(abs_f32_device(cutlass.Float32(row_values[value_idx])))
+                        row_amax = _tree_reduce_max_nan_f32(level)
+                        for shuffle_idx in cutlass.range_constexpr(3):
+                            other_amax = cute.arch.shuffle_sync_bfly(
+                                row_amax,
+                                offset=1 << shuffle_idx,
+                            )
+                            row_amax = max_nan_f32_device(row_amax, other_amax)
+
+                        row_amax = cutlass.max(row_amax, cutlass.Float32(_MIN_AMAX))
+                        sf_fp32 = ceil_to_ue8m0_device(row_amax / cutlass.Float32(_FP8_E4M3_MAX))
+                        sf_bits = cutlass.Int32(
+                            llvm.bitcast(cutlass.Int32.mlir_type, sf_fp32.ir_value())
+                        )
+                        ue8m0_byte = (sf_bits >> cutlass.Int32(23)) & cutlass.Int32(0xFF)
+                        if lane_in_group == 0 and global_m < runtime_m:
+                            stg_u8_raw(
+                                sf_out_addr
+                                + cutlass.Int64(sf_base_offset + global_m * cutlass.Int32(4)),
+                                ue8m0_byte,
+                            )
+
+                        sf_rcp = cutlass.Float32(1.0) / sf_fp32
+                        scaled_values = cute.make_rmem_tensor((16,), cutlass.Float32)
+                        for value_idx in cutlass.range_constexpr(16):
+                            scaled_values[value_idx] = (
+                                cutlass.Float32(row_values[value_idx]) * sf_rcp
+                            )
+                        fp8_values = scaled_values.load().to(self.c_dtype)
+                        fp8_subtile_idx = flat_col_base >> cutlass.Int32(5)
+                        fp8_c_buffer = (
+                            num_prev_subtiles + fp8_subtile_idx + cutlass.Int32(1)
+                        ) % self.num_c_stage
+                        fp8_vector = cute.make_rmem_tensor((16,), self.c_dtype)
+                        fp8_vector.store(fp8_values)
+                        fp8_dest = fp8_thr_copy_r2s.partition_D(sC[(local_row, None, fp8_c_buffer)])
+                        fp8_fragment = cute.make_fragment_like(fp8_dest)
+                        fp8_fragment.store(fp8_vector.load())
+                        cute.copy(
+                            fp8_copy_atom_r2s,
+                            fp8_fragment,
+                            fp8_dest,
+                        )
+
+                    cute.arch.fence_view_async_shared()
                     self.epilog_sync_barrier.arrive_and_wait()
+                    for subtile_idx in cutlass.range(subtile_cnt):
+                        num_prev_subtiles = num_prev_subtiles + 1
+                        c_buffer = num_prev_subtiles % self.num_c_stage
+                        if warp_idx == self.epilog_warp_id[0]:
+                            cute.copy(
+                                tma_atom_c,
+                                bSG_sC[(None, c_buffer)],
+                                bSG_gC[(None, subtile_idx)],
+                            )
+                            c_pipeline.producer_commit()
+                            c_pipeline.producer_acquire()
+                        self.epilog_sync_barrier.arrive_and_wait()
+
+                else:
+                    # Standard BF16 epilogue.
+                    for subtile_idx in cutlass.range(subtile_cnt):
+                        # Load and convert one accumulator subtile.
+                        tTR_tAcc_mn = tTR_tAcc[(None, None, None, subtile_idx)]
+                        cute.copy(tiled_copy_t2r, tTR_tAcc_mn, tTR_rAcc)
+
+                        acc_vec = tiled_copy_r2s.retile(tTR_rAcc).load()
+                        acc_vec = epilogue_op(acc_vec.to(self.c_dtype))
+                        tRS_rC.store(acc_vec)
+
+                        # Stage output for TMA.
+                        num_prev_subtiles = num_prev_subtiles + 1
+                        c_buffer = num_prev_subtiles % self.num_c_stage
+                        cute.copy(
+                            tiled_copy_r2s,
+                            tRS_rC,
+                            tRS_sC[(None, None, None, c_buffer)],
+                        )
+                        # Publish SMEM before TMA.
+                        cute.arch.fence_view_async_shared()
+                        self.epilog_sync_barrier.arrive_and_wait()
+
+                        if warp_idx == self.epilog_warp_id[0]:
+                            cute.copy(
+                                tma_atom_c,
+                                bSG_sC[(None, c_buffer)],
+                                bSG_gC[(None, subtile_idx)],
+                            )
+                            # Commit before reusing this stage.
+                            c_pipeline.producer_commit()
+                            c_pipeline.producer_acquire()
+                        self.epilog_sync_barrier.arrive_and_wait()
 
                 #
                 # Async arrive accumulator buffer empty
                 #
-                epi_pipeline.consumer_release(epi_consumer_state)
-                epi_consumer_state.advance()
+                if cutlass.const_expr(not self.fp8_smem_epi_mode):
+                    epi_pipeline.consumer_release(epi_consumer_state)
+                    epi_consumer_state.advance()
 
                 #
                 # Advance to next tile
@@ -2055,6 +2536,26 @@ def epilog_smem_copy_and_partition(
         tRS_rC = tiled_copy_r2s.retile(tTR_rC)
         return tiled_copy_r2s, tRS_rC, tRS_sC
 
+    def epilog_bf16_smem_copy_and_partition(
+        self,
+        tiled_copy_t2r: cute.TiledCopy,
+        tTR_rC: cute.Tensor,
+        tidx: cutlass.Int32,
+        sC: cute.Tensor,
+    ) -> Tuple[cute.TiledCopy, cute.Tensor, cute.Tensor]:
+        """Partition a BF16 staging tile using the TMEM epilogue ownership."""
+        copy_atom_r2s = sm100_utils.get_smem_store_op(
+            self.c_layout,
+            cutlass.BFloat16,
+            self.acc_dtype,
+            tiled_copy_t2r,
+        )
+        tiled_copy_r2s = cute.make_tiled_copy_D(copy_atom_r2s, tiled_copy_t2r)
+        thr_copy_r2s = tiled_copy_r2s.get_slice(tidx)
+        tRS_sC = thr_copy_r2s.partition_D(sC)
+        tRS_rC = tiled_copy_r2s.retile(tTR_rC)
+        return tiled_copy_r2s, tRS_rC, tRS_sC
+
     def epilog_gmem_copy_and_partition(
         self,
         tidx: cutlass.Int32,
@@ -2116,6 +2617,7 @@ def _compute_stages(
         sfb_count: int,
         num_smem_capacity: int,
         occupancy: int,
+        smem_reserve_bytes: int,
     ) -> Tuple[int, int, int]:
         """Computes the number of stages for A/B/C operands based on heuristics.
 
@@ -2177,8 +2679,6 @@ def _compute_stages(
         ab_bytes_per_stage = cute.size_in_bytes(
             a_dtype, a_smem_layout_stage_one
         ) + cute.size_in_bytes(b_dtype, b_smem_layout_staged_one)
-        # 1024B alignment
-        mbar_helpers_bytes = 1024
         c_bytes_per_stage = cute.size_in_bytes(c_dtype, c_smem_layout_staged_one)
         c_bytes = c_bytes_per_stage * num_c_stage
         sfa_bytes = sfa_count * (sfa_dtype.width // 8) * num_scale_stage
@@ -2190,7 +2690,7 @@ def _compute_stages(
         # Subtract reserved bytes and initial C stages bytes
         # Divide remaining by bytes needed per A/B stage
         num_ab_stage = (
-            num_smem_capacity // occupancy - (mbar_helpers_bytes + c_bytes + scale_bytes)
+            num_smem_capacity // occupancy - (smem_reserve_bytes + c_bytes + scale_bytes)
         ) // ab_bytes_per_stage
 
         # Refine epilogue stages:
@@ -2199,7 +2699,7 @@ def _compute_stages(
         num_c_stage += (
             num_smem_capacity
             - occupancy * ab_bytes_per_stage * num_ab_stage
-            - occupancy * (mbar_helpers_bytes + c_bytes + scale_bytes)
+            - occupancy * (smem_reserve_bytes + c_bytes + scale_bytes)
         ) // (occupancy * c_bytes_per_stage)
         return num_acc_stage, num_ab_stage, num_c_stage, num_scale_stage, num_tile_stage
 
@@ -2293,7 +2793,12 @@ def is_valid_dtypes(
             is_valid = False
         if acc_dtype not in {cutlass.Float32}:
             is_valid = False
-        if c_dtype not in {cutlass.Float32, cutlass.Float16, cutlass.BFloat16}:
+        if c_dtype not in {
+            cutlass.Float32,
+            cutlass.Float16,
+            cutlass.BFloat16,
+            cutlass.Float8E4M3FN,
+        }:
             is_valid = False
         return is_valid
 
@@ -2542,3 +3047,64 @@ def wrapper(
             max_active_clusters,
             stream,
         )
+
+    @cute.jit
+    def wrapper_fp8sf(
+        self,
+        m: cutlass.Int32,
+        n: cutlass.Int32,
+        k: cutlass.Int32,
+        sf_m: cutlass.Int32,
+        sf_n: cutlass.Int32,
+        sf_k: cutlass.Int32,
+        batch_size: cutlass.Int32,
+        a_ptr: cute.Pointer,
+        b_ptr: cute.Pointer,
+        a_sf_ptr: cute.Pointer,
+        b_sf_ptr: cute.Pointer,
+        c_tensor: cute.Tensor,
+        max_active_clusters: cutlass.Constexpr,
+        stream: cuda.CUstream,
+        sf_out_ptr: Optional[cute.Pointer],
+        sf_aligned_mn: cutlass.Int32,
+        n_tiles_per_group: cutlass.Int32,
+        fp8_smem_row_iters: cutlass.Constexpr,
+        use_fp8_smem_epilogue: cutlass.Constexpr,
+    ):
+        """Run blockwise GEMM with fused FP8 and packed-scale output."""
+        a_tensor = cute.make_tensor(
+            a_ptr,
+            layout=cute.make_ordered_layout((m, k, batch_size), order=(1, 0, 2)),
+        )
+        b_tensor = cute.make_tensor(
+            b_ptr,
+            layout=cute.make_ordered_layout((n, k, batch_size), order=(1, 0, 2)),
+        )
+        sfa_tensor = cute.make_tensor(
+            a_sf_ptr,
+            layout=cute.make_ordered_layout((sf_m, sf_k, batch_size), order=(0, 1, 2)),
+        )
+        sfb_tensor = cute.make_tensor(
+            b_sf_ptr,
+            layout=cute.make_ordered_layout((sf_n, sf_k, batch_size), order=(1, 0, 2)),
+        )
+        # Carry the scale address through CuTe's tensor ABI.
+        if cutlass.const_expr(sf_out_ptr is not None):
+            sf_out_tensor_arg = cute.make_tensor(sf_out_ptr, cute.make_layout(1))
+        else:
+            sf_out_tensor_arg = None
+
+        self(
+            a_tensor,
+            b_tensor,
+            c_tensor,
+            sfa_tensor,
+            sfb_tensor,
+            max_active_clusters,
+            stream,
+            sf_out_tensor=sf_out_tensor_arg,
+            sf_aligned_mn=sf_aligned_mn,
+            n_tiles_per_group=n_tiles_per_group,
+            fp8_smem_row_iters=fp8_smem_row_iters,
+            use_fp8_smem_epilogue=use_fp8_smem_epilogue,
+        )
diff --git a/tensorrt_llm/_torch/cute_dsl_kernels/blackwell/dense_blockscaled_gemm_persistent.py b/tensorrt_llm/_torch/cute_dsl_kernels/blackwell/dense_blockscaled_gemm_persistent.py
index 885bebc3dc26..4015888a79e8 100644
--- a/tensorrt_llm/_torch/cute_dsl_kernels/blackwell/dense_blockscaled_gemm_persistent.py
+++ b/tensorrt_llm/_torch/cute_dsl_kernels/blackwell/dense_blockscaled_gemm_persistent.py
@@ -53,11 +53,18 @@
 import cutlass.utils.blackwell_helpers as sm100_utils
 import cutlass.utils.blockscaled_layout as blockscaled_utils
 from cutlass.cute.nvgpu import cpasync, tcgen05
+from cutlass.cutlass_dsl import min as dsl_min
 
+from . import dynamic_mxf8_mainloop
 from .custom_pipeline import PipelineTmaUmma, PipelineUmmaAsync
 from .utils import (TRTLLM_ENABLE_PDL, griddepcontrol_launch_dependents,
                     griddepcontrol_wait, is_power_of_2)
 
+_SCALE_BLOCK_K = 128
+_SCALES_PER_WORD = 4
+_PACKED_SCALE_K = _SCALE_BLOCK_K * _SCALES_PER_WORD
+_PACKED_SCALE_BYTES = 4
+
 
 class Sm100BlockScaledPersistentDenseGemmKernel:
     """Implements batched matrix multiplication (C = A x SFA x B x SFB) with support for various data types
@@ -98,6 +105,12 @@ def __init__(
         mma_tiler_mn: Tuple[int, int],
         cluster_shape_mn: Tuple[int, int],
         use_prefetch: bool = False,
+        packed_k128_scales: bool = False,
+        tma_packed_scales: bool = False,
+        dynamic_mma_n: bool = False,
+        apply_alpha: bool = True,
+        max_ab_stages: Optional[int] = None,
+        l2_swizzle: bool = False,
     ):
         """Initializes the configuration for a Blackwell dense GEMM kernel.
 
@@ -124,6 +137,12 @@ def __init__(
         # K dimension is deferred in _setup_attributes
         self.mma_tiler = (*mma_tiler_mn, 1)
         self.use_prefetch = use_prefetch
+        self.packed_k128_scales = packed_k128_scales
+        self.tma_packed_scales = tma_packed_scales
+        self.max_ab_stages = max_ab_stages
+        self.dynamic_mma_n = dynamic_mma_n
+        self.apply_alpha = apply_alpha
+        self.l2_swizzle = l2_swizzle
         self.cta_group = (tcgen05.CtaGroup.TWO
                           if self.use_2cta_instrs else tcgen05.CtaGroup.ONE)
 
@@ -137,8 +156,15 @@ def __init__(
         )
         self.mma_warp_id = 4
         self.tma_warp_id = 5
-        self.threads_per_cta = 32 * len(
-            (self.mma_warp_id, self.tma_warp_id, *self.epilog_warp_id))
+        self.scale_warp_id = 6
+        specialized_warps = (
+            self.mma_warp_id,
+            self.tma_warp_id,
+            *self.epilog_warp_id,
+        )
+        if self.packed_k128_scales:
+            specialized_warps = (*specialized_warps, self.scale_warp_id)
+        self.threads_per_cta = 32 * len(specialized_warps)
         # Set barrier id for cta sync, epilogue sync and tmem ptr sync
         self.cta_sync_bar_id = 0
         self.epilog_sync_bar_id = 1
@@ -245,6 +271,9 @@ def _setup_attributes(self):
             self.c_layout,
             self.c_dtype,
         )
+        default_epi_tile_n = cute.size(self.epi_tile[1])
+        if self.cta_tile_shape_mnk[1] % default_epi_tile_n != 0:
+            self.epi_tile = (self.epi_tile[0], cute.make_layout(16))
 
         self.epi_tile_n = cute.size(self.epi_tile[1])
 
@@ -264,6 +293,16 @@ def _setup_attributes(self):
             self.smem_capacity,
             self.occupancy,
         )
+        if self.max_ab_stages is not None:
+            self.num_ab_stage = min(self.num_ab_stage, self.max_ab_stages)
+        if self.tma_packed_scales:
+            self.num_c_stage = 2
+        self.num_scale_stage = self.num_ab_stage
+        if self.packed_k128_scales:
+            self.num_scale_stage = max(
+                2,
+                (self.num_ab_stage + _SCALES_PER_WORD - 1) // _SCALES_PER_WORD,
+            )
 
         # Compute A/B/SFA/SFB/C shared memory layout
         self.a_smem_layout_staged = sm100_utils.make_smem_layout_a(
@@ -282,13 +321,13 @@ def _setup_attributes(self):
             tiled_mma,
             self.mma_tiler,
             self.sf_vec_size,
-            self.num_ab_stage,
+            self.num_scale_stage,
         )
         self.sfb_smem_layout_staged = blockscaled_utils.make_smem_layout_sfb(
             tiled_mma,
             self.mma_tiler,
             self.sf_vec_size,
-            self.num_ab_stage,
+            self.num_scale_stage,
         )
         self.c_smem_layout_staged = sm100_utils.make_smem_layout_epi(
             self.c_dtype,
@@ -314,6 +353,121 @@ def _setup_attributes(self):
         # TODO: [alel] Currently set prefetch dist to num_ab_stage, we may have more options for prefetch dist auto tuning
         self.prefetch_dist = self.num_ab_stage
 
+    def _swizzle_work_tile(self, tile_coord: cute.Coord,
+                           mC_mnl: cute.Tensor) -> cute.Coord:
+        if cutlass.const_expr(not self.l2_swizzle):
+            return tile_coord
+
+        cluster_m = self.cluster_shape_mn[0]
+        cluster_n = self.cluster_shape_mn[1]
+        num_ctas_m = cute.ceil_div(cute.size(mC_mnl, mode=[0]),
+                                   self.cta_tile_shape_mnk[0])
+        num_ctas_n = cute.ceil_div(cute.size(mC_mnl, mode=[1]),
+                                   self.cta_tile_shape_mnk[1])
+        num_clusters_m = cute.ceil_div(num_ctas_m, cluster_m)
+        num_clusters_n = cute.ceil_div(num_ctas_n, cluster_n)
+
+        cluster_coord_m = tile_coord[0] // cluster_m
+        cluster_coord_n = tile_coord[1] // cluster_n
+        linear_cluster = cluster_coord_m + num_clusters_m * cluster_coord_n
+
+        # Group clusters like DeepGEMM for L2 reuse.
+        group_ctas = (16 if self.cta_tile_shape_mnk[1] >= 128 else 8)
+        group_clusters = group_ctas // cluster_m
+        clusters_per_group = group_clusters * num_clusters_n
+        group_idx = linear_cluster // clusters_per_group
+        in_group_idx = linear_cluster % clusters_per_group
+        first_cluster_m = group_idx * group_clusters
+        remaining_clusters = num_clusters_m - first_cluster_m
+        valid_group_clusters = dsl_min(cutlass.Int32(group_clusters),
+                                       remaining_clusters)
+
+        mapped_cluster_m = (first_cluster_m +
+                            in_group_idx % valid_group_clusters)
+        mapped_cluster_n = in_group_idx // valid_group_clusters
+        return (
+            mapped_cluster_m * cluster_m + tile_coord[0] % cluster_m,
+            mapped_cluster_n * cluster_n + tile_coord[1] % cluster_n,
+            tile_coord[2],
+        )
+
+    @cute.jit(preprocess=True)
+    def wrapper_dsv4_splitk_packed_ue8m0(
+        self,
+        m: cutlass.Constexpr,
+        n: cutlass.Constexpr,
+        k: cutlass.Constexpr,
+        num_splits: cutlass.Constexpr,
+        sfa_k_stride: cutlass.Constexpr,
+        sfb_k_stride: cutlass.Constexpr,
+        a_ptr: cute.Pointer,
+        b_ptr: cute.Pointer,
+        sfa_ptr: cute.Pointer,
+        sfb_ptr: cute.Pointer,
+        c_tensor: cute.Tensor,
+        alpha: cute.Tensor,
+        max_active_clusters: cutlass.Constexpr,
+        stream: cuda.CUstream,
+        swap_ab: cutlass.Constexpr = True,
+    ):
+        """Run the DSV4 O_b GEMM and emit split-major partials."""
+        split_k = k // num_splits
+        packed_k_per_split = split_k // _PACKED_SCALE_K
+        a_tensor = cute.make_tensor(
+            a_ptr,
+            cute.make_layout(
+                (m, split_k, num_splits),
+                stride=(k, 1, split_k),
+            ),
+        )
+        b_tensor = cute.make_tensor(
+            b_ptr,
+            cute.make_layout(
+                (n, split_k, num_splits),
+                stride=(k, 1, split_k),
+            ),
+        )
+        sfa_tensor = cute.make_tensor(
+            sfa_ptr,
+            cute.make_layout(
+                (m, packed_k_per_split, num_splits),
+                stride=(
+                    1,
+                    sfa_k_stride,
+                    sfa_k_stride * packed_k_per_split,
+                ),
+            ),
+        )
+        sfb_tensor = cute.make_tensor(
+            sfb_ptr,
+            cute.make_layout(
+                (n, packed_k_per_split, num_splits),
+                stride=(
+                    1,
+                    sfb_k_stride,
+                    sfb_k_stride * packed_k_per_split,
+                ),
+            ),
+        )
+        output_stride = ((1, m, m * n) if swap_ab else (n, 1, m * n))
+        output_tensor = cute.make_tensor(
+            c_tensor.iterator,
+            cute.make_layout(
+                (m, n, num_splits),
+                stride=output_stride,
+            ),
+        )
+        self(
+            a_tensor,
+            b_tensor,
+            sfa_tensor,
+            sfb_tensor,
+            output_tensor,
+            alpha,
+            max_active_clusters,
+            stream,
+        )
+
     @cute.jit
     def __call__(
         self,
@@ -350,7 +504,9 @@ def __call__(
         # Setup static attributes before smem/grid/tma computation
         self.a_dtype: Type[cutlass.Numeric] = a_tensor.element_type
         self.b_dtype: Type[cutlass.Numeric] = b_tensor.element_type
-        self.sf_dtype: Type[cutlass.Numeric] = sfa_tensor.element_type
+        self.sf_dtype: Type[cutlass.Numeric] = (cutlass.Float8E8M0FNU
+                                                if self.packed_k128_scales else
+                                                sfa_tensor.element_type)
         self.c_dtype: Type[cutlass.Numeric] = c_tensor.element_type
         self.a_major_mode = utils.LayoutEnum.from_tensor(
             a_tensor).mma_major_mode()
@@ -366,16 +522,17 @@ def __call__(
         # Setup attributes that dependent on gemm inputs
         self._setup_attributes()
 
-        # Setup sfa/sfb tensor by filling A/B tensor to scale factor atom layout
-        # ((Atom_M, Rest_M),(Atom_K, Rest_K),RestL)
-        sfa_layout = blockscaled_utils.tile_atom_to_shape_SF(
-            a_tensor.shape, self.sf_vec_size)
-        sfa_tensor = cute.make_tensor(sfa_tensor.iterator, sfa_layout)
+        # Packed scales already use their runtime MN-major layout.
+        if cutlass.const_expr(not self.packed_k128_scales):
+            # ((Atom_M, Rest_M),(Atom_K, Rest_K),RestL)
+            sfa_layout = blockscaled_utils.tile_atom_to_shape_SF(
+                a_tensor.shape, self.sf_vec_size)
+            sfa_tensor = cute.make_tensor(sfa_tensor.iterator, sfa_layout)
 
-        # ((Atom_N, Rest_N),(Atom_K, Rest_K),RestL)
-        sfb_layout = blockscaled_utils.tile_atom_to_shape_SF(
-            b_tensor.shape, self.sf_vec_size)
-        sfb_tensor = cute.make_tensor(sfb_tensor.iterator, sfb_layout)
+            # ((Atom_N, Rest_N),(Atom_K, Rest_K),RestL)
+            sfb_layout = blockscaled_utils.tile_atom_to_shape_SF(
+                b_tensor.shape, self.sf_vec_size)
+            sfb_tensor = cute.make_tensor(sfb_tensor.iterator, sfb_layout)
 
         tiled_mma = sm100_utils.make_blockscaled_trivial_tiled_mma(
             self.a_dtype,
@@ -426,56 +583,97 @@ def __call__(
             self.cluster_layout_vmnk.shape,
         )
 
-        # Setup TMA load for SFA
-        sfa_op = sm100_utils.cluster_shape_to_tma_atom_A(
-            self.cluster_shape_mn, tiled_mma.thr_id)
         sfa_smem_layout = cute.slice_(self.sfa_smem_layout_staged,
                                       (None, None, None, 0))
-        tma_atom_sfa, tma_tensor_sfa = cute.nvgpu.make_tiled_tma_atom_A(
-            sfa_op,
-            sfa_tensor,
-            sfa_smem_layout,
-            self.mma_tiler,
-            tiled_mma,
-            self.cluster_layout_vmnk.shape,
-            internal_type=cutlass.Int16,
-        )
-
-        # Setup TMA load for SFB
-        sfb_op = sm100_utils.cluster_shape_to_tma_atom_SFB(
-            self.cluster_shape_mn, tiled_mma.thr_id)
         sfb_smem_layout = cute.slice_(self.sfb_smem_layout_staged,
                                       (None, None, None, 0))
-        tma_atom_sfb, tma_tensor_sfb = cute.nvgpu.make_tiled_tma_atom_B(
-            sfb_op,
-            sfb_tensor,
-            sfb_smem_layout,
-            self.mma_tiler_sfb,
-            tiled_mma_sfb,
-            self.cluster_layout_sfb_vmnk.shape,
-            internal_type=cutlass.Int16,
-        )
-        if cutlass.const_expr(self.cta_tile_shape_mnk[1] == 192):
-            x = tma_tensor_sfb.stride[0][1]
-            y = cute.ceil_div(tma_tensor_sfb.shape[0][1], 4)
-
-            new_shape = ((tma_tensor_sfb.shape[0][0], ((2, 2), y)),
-                         tma_tensor_sfb.shape[1], tma_tensor_sfb.shape[2])
-            # Use right multiplication for ScaledBasis (3 * x instead of x * 3)
-            x_times_3 = 3 * x
-            new_stride = ((tma_tensor_sfb.stride[0][0], ((x, x), x_times_3)),
-                          tma_tensor_sfb.stride[1], tma_tensor_sfb.stride[2])
-            tma_tensor_sfb_new_layout = cute.make_layout(new_shape,
-                                                         stride=new_stride)
-            tma_tensor_sfb = cute.make_tensor(tma_tensor_sfb.iterator,
-                                              tma_tensor_sfb_new_layout)
+        tma_atom_sfa = None
+        tma_tensor_sfa = sfa_tensor
+        tma_atom_sfb = None
+        tma_tensor_sfb = sfb_tensor
+        if cutlass.const_expr(self.packed_k128_scales
+                              and self.tma_packed_scales):
+            # TMA one packed scale word per output row.
+            scale_tma_op = cpasync.CopyBulkTensorTileG2SOp()
+            sfa_raw_stage_layout = cute.make_layout(
+                (self.cta_tile_shape_mnk[0], ), stride=(1, ))
+            tma_atom_sfa, tma_tensor_sfa = cpasync.make_tiled_tma_atom(
+                scale_tma_op,
+                sfa_tensor,
+                sfa_raw_stage_layout,
+                (self.cta_tile_shape_mnk[0], ),
+            )
+            num_sfb_tma_rows = self.cta_tile_shape_mnk[1]
+            sfb_raw_stage_layout = cute.make_layout((num_sfb_tma_rows, ),
+                                                    stride=(1, ))
+            tma_atom_sfb, tma_tensor_sfb = cpasync.make_tiled_tma_atom(
+                scale_tma_op,
+                sfb_tensor,
+                sfb_raw_stage_layout,
+                (num_sfb_tma_rows, ),
+            )
+        elif cutlass.const_expr(not self.packed_k128_scales):
+            # Setup TMA load for SFA
+            sfa_op = sm100_utils.cluster_shape_to_tma_atom_A(
+                self.cluster_shape_mn, tiled_mma.thr_id)
+            tma_atom_sfa, tma_tensor_sfa = cute.nvgpu.make_tiled_tma_atom_A(
+                sfa_op,
+                sfa_tensor,
+                sfa_smem_layout,
+                self.mma_tiler,
+                tiled_mma,
+                self.cluster_layout_vmnk.shape,
+                internal_type=cutlass.Int16,
+            )
+
+            # Setup TMA load for SFB
+            sfb_op = sm100_utils.cluster_shape_to_tma_atom_SFB(
+                self.cluster_shape_mn, tiled_mma.thr_id)
+            tma_atom_sfb, tma_tensor_sfb = cute.nvgpu.make_tiled_tma_atom_B(
+                sfb_op,
+                sfb_tensor,
+                sfb_smem_layout,
+                self.mma_tiler_sfb,
+                tiled_mma_sfb,
+                self.cluster_layout_sfb_vmnk.shape,
+                internal_type=cutlass.Int16,
+            )
+            if cutlass.const_expr(self.cta_tile_shape_mnk[1] == 192):
+                x = tma_tensor_sfb.stride[0][1]
+                y = cute.ceil_div(tma_tensor_sfb.shape[0][1], 4)
+
+                new_shape = ((tma_tensor_sfb.shape[0][0], ((2, 2), y)),
+                             tma_tensor_sfb.shape[1], tma_tensor_sfb.shape[2])
+                # Use right multiplication for ScaledBasis (3 * x instead of x * 3)
+                x_times_3 = 3 * x
+                new_stride = ((tma_tensor_sfb.stride[0][0],
+                               ((x, x), x_times_3)), tma_tensor_sfb.stride[1],
+                              tma_tensor_sfb.stride[2])
+                tma_tensor_sfb_new_layout = cute.make_layout(new_shape,
+                                                             stride=new_stride)
+                tma_tensor_sfb = cute.make_tensor(tma_tensor_sfb.iterator,
+                                                  tma_tensor_sfb_new_layout)
 
         a_copy_size = cute.size_in_bytes(self.a_dtype, a_smem_layout)
         b_copy_size = cute.size_in_bytes(self.b_dtype, b_smem_layout)
-        sfa_copy_size = cute.size_in_bytes(self.sf_dtype, sfa_smem_layout)
-        sfb_copy_size = cute.size_in_bytes(self.sf_dtype, sfb_smem_layout)
-        self.num_tma_load_bytes = (a_copy_size + b_copy_size + sfa_copy_size +
-                                   sfb_copy_size) * atom_thr_size
+        if cutlass.const_expr(self.packed_k128_scales):
+            sfa_copy_size = 0
+            sfb_copy_size = 0
+        else:
+            sfa_copy_size = cute.size_in_bytes(self.sf_dtype, sfa_smem_layout)
+            sfb_copy_size = cute.size_in_bytes(self.sf_dtype, sfb_smem_layout)
+        if cutlass.const_expr(self.packed_k128_scales):
+            self.num_tma_load_bytes = (a_copy_size +
+                                       b_copy_size) * atom_thr_size
+        else:
+            self.num_tma_load_bytes = (a_copy_size + b_copy_size + sfa_copy_size
+                                       + sfb_copy_size) * atom_thr_size
+        self.num_scale_tma_load_bytes = 0
+        if cutlass.const_expr(self.packed_k128_scales
+                              and self.tma_packed_scales):
+            self.num_scale_tma_load_bytes = (
+                self.cta_tile_shape_mnk[0] +
+                self.cta_tile_shape_mnk[1]) * _PACKED_SCALE_BYTES
 
         # Setup TMA store for C
         epi_smem_layout = cute.slice_(self.c_smem_layout_staged,
@@ -496,6 +694,15 @@ def __call__(
         )
 
         self.buffer_align_bytes = 1024
+        sfb_storage_elems = cute.cosize(self.sfb_smem_layout_staged)
+        if cutlass.const_expr(self.packed_k128_scales):
+            sfb_storage_elems = (
+                cute.round_up(self.cta_tile_shape_mnk[1], _SCALE_BLOCK_K) *
+                _PACKED_SCALE_BYTES * self.num_scale_stage)
+        scale_tma_mbar_elems = 0
+        if cutlass.const_expr(self.packed_k128_scales
+                              and self.tma_packed_scales):
+            scale_tma_mbar_elems = self.num_scale_stage * 2
 
         # Define shared storage for kernel
         @cute.struct
@@ -504,6 +711,10 @@ class SharedStorage:
                                                    self.num_ab_stage]
             ab_empty_mbar_ptr: cute.struct.MemRange[cutlass.Int64,
                                                     self.num_ab_stage]
+            scale_mbar_ptr: cute.struct.MemRange[cutlass.Int64,
+                                                 self.num_scale_stage * 2]
+            scale_tma_mbar_ptr: cute.struct.MemRange[cutlass.Int64,
+                                                     scale_tma_mbar_elems]
             acc_full_mbar_ptr: cute.struct.MemRange[cutlass.Int64,
                                                     self.num_acc_stage]
             acc_empty_mbar_ptr: cute.struct.MemRange[cutlass.Int64,
@@ -540,8 +751,7 @@ class SharedStorage:
             ]
             # (MMA, MMA_N, MMA_K, STAGE)
             sSFB: cute.struct.Align[
-                cute.struct.MemRange[self.sf_dtype,
-                                     cute.cosize(self.sfb_smem_layout_staged)],
+                cute.struct.MemRange[self.sf_dtype, sfb_storage_elems],
                 self.buffer_align_bytes,
             ]
 
@@ -593,9 +803,9 @@ def kernel(
         mA_mkl: cute.Tensor,
         tma_atom_b: cute.CopyAtom,
         mB_nkl: cute.Tensor,
-        tma_atom_sfa: cute.CopyAtom,
+        tma_atom_sfa: Optional[cute.CopyAtom],
         mSFA_mkl: cute.Tensor,
-        tma_atom_sfb: cute.CopyAtom,
+        tma_atom_sfb: Optional[cute.CopyAtom],
         mSFB_nkl: cute.Tensor,
         tma_atom_c: Optional[cute.CopyAtom],
         mC_mnl: cute.Tensor,
@@ -614,7 +824,8 @@ def kernel(
         """
         GPU device kernel performing the Persistent batched GEMM computation.
         """
-        alpha_value = alpha[0].to(self.c_dtype)
+        if cutlass.const_expr(self.apply_alpha):
+            alpha_value = alpha[0].to(self.c_dtype)
 
         warp_idx = cute.arch.warp_idx()
         warp_idx = cute.arch.make_warp_uniform(warp_idx)
@@ -625,8 +836,10 @@ def kernel(
         if warp_idx == self.tma_warp_id:
             cpasync.prefetch_descriptor(tma_atom_a)
             cpasync.prefetch_descriptor(tma_atom_b)
-            cpasync.prefetch_descriptor(tma_atom_sfa)
-            cpasync.prefetch_descriptor(tma_atom_sfb)
+            if cutlass.const_expr(not self.packed_k128_scales
+                                  or self.tma_packed_scales):
+                cpasync.prefetch_descriptor(tma_atom_sfa)
+                cpasync.prefetch_descriptor(tma_atom_sfb)
             cpasync.prefetch_descriptor(tma_atom_c)
 
         use_2cta_instrs = cute.size(tiled_mma.thr_id.shape) == 2
@@ -671,11 +884,37 @@ def kernel(
             cta_layout_vmnk=cluster_layout_vmnk,
         )
 
+        if cutlass.const_expr(self.packed_k128_scales):
+            scale_pipeline = pipeline.PipelineAsync.create(
+                barrier_storage=storage.scale_mbar_ptr.data_ptr(),
+                num_stages=self.num_scale_stage,
+                producer_group=pipeline.CooperativeGroup(
+                    pipeline.Agent.Thread,
+                    32 * cute.size(tiled_mma.thr_id.shape),
+                ),
+                consumer_group=pipeline.CooperativeGroup(pipeline.Agent.Thread),
+                producer_mask=cta_rank_in_cluster ^ mma_tile_coord_v,
+                defer_sync=True,
+            )
+            if cutlass.const_expr(self.tma_packed_scales):
+                scale_tma_pipeline = pipeline.PipelineTmaAsync.create(
+                    barrier_storage=storage.scale_tma_mbar_ptr.data_ptr(),
+                    num_stages=self.num_scale_stage,
+                    producer_group=pipeline.CooperativeGroup(
+                        pipeline.Agent.Thread),
+                    consumer_group=pipeline.CooperativeGroup(
+                        pipeline.Agent.Thread),
+                    tx_count=self.num_scale_tma_load_bytes,
+                    cta_layout_vmnk=cute.make_layout((1, 1, 1, 1)),
+                    tidx=tidx,
+                    defer_sync=True,
+                )
+
         # Initialize acc_pipeline (barrier) and states
         acc_pipeline_producer_group = pipeline.CooperativeGroup(
             pipeline.Agent.Thread)
         num_acc_consumer_threads = len(
-            self.epilog_warp_id) * (2 if use_2cta_instrs else 1)
+            self.epilog_warp_id) * (2 if self.use_2cta_instrs else 1)
         acc_pipeline_consumer_group = pipeline.CooperativeGroup(
             pipeline.Agent.Thread, num_acc_consumer_threads)
         acc_pipeline = PipelineUmmaAsync.create(
@@ -687,7 +926,7 @@ def kernel(
         )
 
         # Tensor memory dealloc barrier init
-        if use_2cta_instrs:
+        if cutlass.const_expr(self.use_2cta_instrs):
             if warp_idx == self.tma_warp_id:
                 num_tmem_dealloc_threads = 32
                 with cute.arch.elect_one():
@@ -715,6 +954,40 @@ def kernel(
         sSFA = storage.sSFA.get_tensor(sfa_smem_layout_staged)
         # (MMA, MMA_N, MMA_K, STAGE)
         sSFB = storage.sSFB.get_tensor(sfb_smem_layout_staged)
+        if cutlass.const_expr(self.packed_k128_scales):
+            num_sfb_rows = cute.round_up(self.cta_tile_shape_mnk[1],
+                                         _SCALE_BLOCK_K)
+            sSFA_raw = cute.make_tensor(
+                cute.recast_ptr(storage.sSFA.data_ptr(), dtype=cutlass.Uint32),
+                cute.make_layout(
+                    (self.cta_tile_shape_mnk[0], self.num_scale_stage),
+                    stride=(1, self.cta_tile_shape_mnk[0]),
+                ),
+            )
+            sSFB = cute.make_tensor(
+                sSFB.iterator,
+                cute.make_layout(
+                    sSFB.shape,
+                    stride=(sSFB.stride[0], sSFB.stride[1], sSFB.stride[2],
+                            num_sfb_rows * _PACKED_SCALE_BYTES),
+                ),
+            )
+            sSFB_raw = cute.make_tensor(
+                cute.recast_ptr(storage.sSFB.data_ptr(), dtype=cutlass.Uint32),
+                cute.make_layout(
+                    (num_sfb_rows, self.num_scale_stage),
+                    stride=(1, num_sfb_rows),
+                ),
+            )
+            if cutlass.const_expr(self.tma_packed_scales):
+                sSFB_tma_view = cute.make_tensor(
+                    cute.recast_ptr(storage.sSFB.data_ptr(),
+                                    dtype=cutlass.Uint32),
+                    cute.make_layout(
+                        (self.cta_tile_shape_mnk[1], self.num_scale_stage),
+                        stride=(1, num_sfb_rows),
+                    ),
+                )
 
         #
         # Compute multicast mask for A/B/SFA/SFB buffer full
@@ -747,14 +1020,30 @@ def kernel(
         gB_nkl = cute.local_tile(mB_nkl,
                                  cute.slice_(self.mma_tiler, (0, None, None)),
                                  (None, None, None))
-        # (bM, bK, RestM, RestK, RestL)
-        gSFA_mkl = cute.local_tile(mSFA_mkl,
-                                   cute.slice_(self.mma_tiler, (None, 0, None)),
-                                   (None, None, None))
-        # (bN, bK, RestN, RestK, RestL)
-        gSFB_nkl = cute.local_tile(
-            mSFB_nkl, cute.slice_(self.mma_tiler_sfb, (0, None, None)),
-            (None, None, None))
+        if cutlass.const_expr(not self.packed_k128_scales):
+            # (bM, bK, RestM, RestK, RestL)
+            gSFA_mkl = cute.local_tile(
+                mSFA_mkl,
+                cute.slice_(self.mma_tiler, (None, 0, None)),
+                (None, None, None),
+            )
+            # (bN, bK, RestN, RestK, RestL)
+            gSFB_nkl = cute.local_tile(
+                mSFB_nkl,
+                cute.slice_(self.mma_tiler_sfb, (0, None, None)),
+                (None, None, None),
+            )
+        elif cutlass.const_expr(self.tma_packed_scales):
+            gSFA_tma_mkl = cute.local_tile(
+                mSFA_mkl,
+                (self.cta_tile_shape_mnk[0], ),
+                (None, None, None),
+            )
+            gSFB_tma_nkl = cute.local_tile(
+                mSFB_nkl,
+                (self.cta_tile_shape_mnk[1], ),
+                (None, None, None),
+            )
         # (bM, bN, RestM, RestN, RestL)
         gC_mnl = cute.local_tile(mC_mnl,
                                  cute.slice_(self.mma_tiler, (None, None, 0)),
@@ -770,10 +1059,11 @@ def kernel(
         tCgA = thr_mma.partition_A(gA_mkl)
         # (MMA, MMA_N, MMA_K, RestN, RestK, RestL)
         tCgB = thr_mma.partition_B(gB_nkl)
-        # (MMA, MMA_M, MMA_K, RestM, RestK, RestL)
-        tCgSFA = thr_mma.partition_A(gSFA_mkl)
-        # (MMA, MMA_N, MMA_K, RestN, RestK, RestL)
-        tCgSFB = thr_mma_sfb.partition_B(gSFB_nkl)
+        if cutlass.const_expr(not self.packed_k128_scales):
+            # (MMA, MMA_M, MMA_K, RestM, RestK, RestL)
+            tCgSFA = thr_mma.partition_A(gSFA_mkl)
+            # (MMA, MMA_N, MMA_K, RestN, RestK, RestL)
+            tCgSFB = thr_mma_sfb.partition_B(gSFB_nkl)
         # (MMA, MMA_M, MMA_N, RestM, RestN, RestL)
         tCgC = thr_mma.partition_C(gC_mnl)
 
@@ -805,34 +1095,48 @@ def kernel(
             cute.group_modes(tCgB, 0, 3),
         )
 
-        #  TMA load SFA partition_S/D
-        sfa_cta_layout = a_cta_layout
-        # ((atom_v, rest_v), STAGE)
-        # ((atom_v, rest_v), RestM, RestK, RestL)
-        tAsSFA, tAgSFA = cute.nvgpu.cpasync.tma_partition(
-            tma_atom_sfa,
-            block_in_cluster_coord_vmnk[2],
-            sfa_cta_layout,
-            cute.group_modes(sSFA, 0, 3),
-            cute.group_modes(tCgSFA, 0, 3),
-        )
-        tAsSFA = cute.filter_zeros(tAsSFA)
-        tAgSFA = cute.filter_zeros(tAgSFA)
-
-        # TMA load SFB partition_S/D
-        sfb_cta_layout = cute.make_layout(
-            cute.slice_(cluster_layout_sfb_vmnk, (0, None, 0, 0)).shape)
-        # ((atom_v, rest_v), STAGE)
-        # ((atom_v, rest_v), RestN, RestK, RestL)
-        tBsSFB, tBgSFB = cute.nvgpu.cpasync.tma_partition(
-            tma_atom_sfb,
-            block_in_cluster_coord_sfb_vmnk[1],
-            sfb_cta_layout,
-            cute.group_modes(sSFB, 0, 3),
-            cute.group_modes(tCgSFB, 0, 3),
-        )
-        tBsSFB = cute.filter_zeros(tBsSFB)
-        tBgSFB = cute.filter_zeros(tBgSFB)
+        if cutlass.const_expr(not self.packed_k128_scales):
+            # TMA load SFA partition_S/D
+            sfa_cta_layout = a_cta_layout
+            tAsSFA, tAgSFA = cute.nvgpu.cpasync.tma_partition(
+                tma_atom_sfa,
+                block_in_cluster_coord_vmnk[2],
+                sfa_cta_layout,
+                cute.group_modes(sSFA, 0, 3),
+                cute.group_modes(tCgSFA, 0, 3),
+            )
+            tAsSFA = cute.filter_zeros(tAsSFA)
+            tAgSFA = cute.filter_zeros(tAgSFA)
+
+            # TMA load SFB partition_S/D
+            sfb_cta_layout = cute.make_layout(
+                cute.slice_(cluster_layout_sfb_vmnk, (0, None, 0, 0)).shape)
+            # ((atom_v, rest_v), STAGE)
+            # ((atom_v, rest_v), RestN, RestK, RestL)
+            tBsSFB, tBgSFB = cute.nvgpu.cpasync.tma_partition(
+                tma_atom_sfb,
+                block_in_cluster_coord_sfb_vmnk[1],
+                sfb_cta_layout,
+                cute.group_modes(sSFB, 0, 3),
+                cute.group_modes(tCgSFB, 0, 3),
+            )
+            tBsSFB = cute.filter_zeros(tBsSFB)
+            tBgSFB = cute.filter_zeros(tBgSFB)
+        elif cutlass.const_expr(self.tma_packed_scales):
+            tAsSFA_tma, tAgSFA_tma = cpasync.tma_partition(
+                tma_atom_sfa,
+                0,
+                cute.make_layout(1),
+                sSFA_raw,
+                gSFA_tma_mkl,
+            )
+            tBsSFB_tma, tBgSFB_tma = cpasync.tma_partition(
+                tma_atom_sfb,
+                0,
+                cute.make_layout(1),
+                sSFB_tma_view,
+                gSFB_tma_nkl,
+            )
 
         #
         # Partition shared/tensor memory tensor for TiledMMA_A/B/C
@@ -886,16 +1190,22 @@ def kernel(
 
             ab_producer_state = pipeline.make_pipeline_state(
                 pipeline.PipelineUserType.Producer, self.num_ab_stage)
-
             while work_tile.is_valid_tile:
 
                 # Get tile coord from tile scheduler
-                cur_tile_coord = work_tile.tile_idx
+                cur_tile_coord = self._swizzle_work_tile(
+                    work_tile.tile_idx, mC_mnl)
                 mma_tile_coord_mnl = (
                     cur_tile_coord[0] // cute.size(tiled_mma.thr_id.shape),
                     cur_tile_coord[1],
                     cur_tile_coord[2],
                 )
+                valid_mma_n = cutlass.Int32(self.mma_tiler[1])
+                if cutlass.const_expr(self.dynamic_mma_n):
+                    remaining_n = (cute.size(mB_nkl, mode=[0]) -
+                                   mma_tile_coord_mnl[1] * self.mma_tiler[1])
+                    if remaining_n < valid_mma_n:
+                        valid_mma_n = cutlass.Int32(remaining_n)
 
                 #
                 # Slice to per mma tile index
@@ -904,18 +1214,77 @@ def kernel(
                 tAgA_slice = tAgA[(None, mma_tile_coord_mnl[0], None,
                                    mma_tile_coord_mnl[2])]
                 # ((atom_v, rest_v), RestK)
-                tBgB_slice = tBgB[(None, mma_tile_coord_mnl[1], None,
-                                   mma_tile_coord_mnl[2])]
-
-                # ((atom_v, rest_v), RestK)
-                tAgSFA_slice = tAgSFA[(None, mma_tile_coord_mnl[0], None,
-                                       mma_tile_coord_mnl[2])]
-                slice_n = mma_tile_coord_mnl[1]
-                if cutlass.const_expr(self.cta_tile_shape_mnk[1] == 64):
-                    slice_n = mma_tile_coord_mnl[1] // 2
-                # ((atom_v, rest_v), RestK)
-                tBgSFB_slice = tBgSFB[(None, slice_n, None,
+                if cutlass.const_expr(self.dynamic_mma_n):
+                    real_b = mB_nkl
+                    if cutlass.const_expr(self.use_2cta_instrs):
+                        load_shift = cutlass.Int32(0)
+                        if not is_leader_cta:
+                            load_shift = dynamic_mxf8_mainloop.compute_non_leader_cta_load_shift(
+                                valid_n=valid_mma_n,
+                                mma_tiler_n=self.mma_tiler[1],
+                            )
+                        real_b = cute.domain_offset((load_shift, 0, 0), real_b)
+                    gB_dynamic = cute.local_tile(
+                        real_b,
+                        cute.slice_(self.mma_tiler, (0, None, None)),
+                        (None, None, None),
+                    )
+                    tCgB_dynamic = thr_mma.partition_B(gB_dynamic)
+                    tBsB_dynamic, tBgB_dynamic = cpasync.tma_partition(
+                        tma_atom_b,
+                        block_in_cluster_coord_vmnk[1],
+                        b_cta_layout,
+                        cute.group_modes(sB, 0, 3),
+                        cute.group_modes(tCgB_dynamic, 0, 3),
+                    )
+                    tBgB_slice = tBgB_dynamic[(None, mma_tile_coord_mnl[1],
+                                               None, mma_tile_coord_mnl[2])]
+                    tBsB_for_load = tBsB_dynamic
+                else:
+                    tBgB_slice = tBgB[(None, mma_tile_coord_mnl[1], None,
                                        mma_tile_coord_mnl[2])]
+                    tBsB_for_load = tBsB
+
+                if cutlass.const_expr(not self.packed_k128_scales):
+                    # ((atom_v, rest_v), RestK)
+                    tAgSFA_slice = tAgSFA[(None, mma_tile_coord_mnl[0], None,
+                                           mma_tile_coord_mnl[2])]
+                    if cutlass.const_expr(self.dynamic_mma_n):
+                        sfb_load_offset = (mma_tile_coord_mnl[1] *
+                                           self.mma_tiler[1])
+                        sfb_offset_coord = cute.idx2crd(sfb_load_offset,
+                                                        mSFB_nkl.shape[0])
+                        sfb_k_zero = cute.idx2crd(0, mSFB_nkl.shape[1])
+                        sfb_l_zero = cute.idx2crd(0, mSFB_nkl.shape[2])
+                        real_sfb = cute.domain_offset(
+                            (sfb_offset_coord, sfb_k_zero, sfb_l_zero),
+                            mSFB_nkl)
+                        gSFB_dynamic = cute.local_tile(
+                            real_sfb,
+                            cute.slice_(self.mma_tiler_sfb, (0, None, None)),
+                            (None, None, None),
+                        )
+                        tCgSFB_dynamic = thr_mma_sfb.partition_B(gSFB_dynamic)
+                        tBsSFB_dynamic, tBgSFB_dynamic = (
+                            cute.nvgpu.cpasync.tma_partition(
+                                tma_atom_sfb,
+                                block_in_cluster_coord_sfb_vmnk[1],
+                                sfb_cta_layout,
+                                cute.group_modes(sSFB, 0, 3),
+                                cute.group_modes(tCgSFB_dynamic, 0, 3),
+                            ))
+                        tBsSFB_for_load = cute.filter_zeros(tBsSFB_dynamic)
+                        tBgSFB_dynamic = cute.filter_zeros(tBgSFB_dynamic)
+                        tBgSFB_slice = tBgSFB_dynamic[(None, 0, None,
+                                                       mma_tile_coord_mnl[2])]
+                    else:
+                        slice_n = mma_tile_coord_mnl[1]
+                        if cutlass.const_expr(self.cta_tile_shape_mnk[1] == 64):
+                            slice_n = mma_tile_coord_mnl[1] // 2
+                        # ((atom_v, rest_v), RestK)
+                        tBgSFB_slice = tBgSFB[(None, slice_n, None,
+                                               mma_tile_coord_mnl[2])]
+                        tBsSFB_for_load = tBsSFB
 
                 if cutlass.const_expr(self.use_prefetch):
                     # Prefetch both A and B (default behavior)
@@ -931,14 +1300,15 @@ def kernel(
                             tma_atom_b,
                             tBgB_slice[(None, pf_k_block)],
                         )
-                        cute.prefetch(
-                            tma_atom_sfa,
-                            tAgSFA_slice[(None, pf_k_block)],
-                        )
-                        cute.prefetch(
-                            tma_atom_sfb,
-                            tBgSFB_slice[(None, pf_k_block)],
-                        )
+                        if cutlass.const_expr(not self.packed_k128_scales):
+                            cute.prefetch(
+                                tma_atom_sfa,
+                                tAgSFA_slice[(None, pf_k_block)],
+                            )
+                            cute.prefetch(
+                                tma_atom_sfb,
+                                tBgSFB_slice[(None, pf_k_block)],
+                            )
 
                 # Peek (try_wait) AB buffer empty for k_block = prefetch_k_block_cnt
                 ab_producer_state.reset_count()
@@ -950,11 +1320,29 @@ def kernel(
                 # Tma load loop
                 #
                 for k_block in cutlass.range(0, k_block_cnt, 1, unroll=1):
-                    # Conditionally wait for AB buffer empty
                     ab_pipeline.producer_acquire(ab_producer_state,
                                                  peek_ab_empty_status)
 
                     # TMA load A/B/SFA/SFB
+                    sf_k_block = ab_producer_state.count
+                    if cutlass.const_expr(not self.packed_k128_scales):
+                        cute.copy(
+                            tma_atom_sfa,
+                            tAgSFA_slice[(None, sf_k_block)],
+                            tAsSFA[(None, ab_producer_state.index)],
+                            tma_bar_ptr=ab_pipeline.producer_get_barrier(
+                                ab_producer_state),
+                            mcast_mask=sfa_full_mcast_mask,
+                        )
+                        cute.copy(
+                            tma_atom_sfb,
+                            tBgSFB_slice[(None, sf_k_block)],
+                            tBsSFB_for_load[(None, ab_producer_state.index)],
+                            tma_bar_ptr=ab_pipeline.producer_get_barrier(
+                                ab_producer_state),
+                            mcast_mask=sfb_full_mcast_mask,
+                        )
+
                     cute.copy(
                         tma_atom_a,
                         tAgA_slice[(None, ab_producer_state.count)],
@@ -966,27 +1354,11 @@ def kernel(
                     cute.copy(
                         tma_atom_b,
                         tBgB_slice[(None, ab_producer_state.count)],
-                        tBsB[(None, ab_producer_state.index)],
+                        tBsB_for_load[(None, ab_producer_state.index)],
                         tma_bar_ptr=ab_pipeline.producer_get_barrier(
                             ab_producer_state),
                         mcast_mask=b_full_mcast_mask,
                     )
-                    cute.copy(
-                        tma_atom_sfa,
-                        tAgSFA_slice[(None, ab_producer_state.count)],
-                        tAsSFA[(None, ab_producer_state.index)],
-                        tma_bar_ptr=ab_pipeline.producer_get_barrier(
-                            ab_producer_state),
-                        mcast_mask=sfa_full_mcast_mask,
-                    )
-                    cute.copy(
-                        tma_atom_sfb,
-                        tBgSFB_slice[(None, ab_producer_state.count)],
-                        tBsSFB[(None, ab_producer_state.index)],
-                        tma_bar_ptr=ab_pipeline.producer_get_barrier(
-                            ab_producer_state),
-                        mcast_mask=sfb_full_mcast_mask,
-                    )
 
                     # Prefetch: Rolling prefetch for next tiles
                     if cutlass.const_expr(self.use_prefetch):
@@ -1001,14 +1373,15 @@ def kernel(
                                 tma_atom_b,
                                 tBgB_slice[(None, future_k_block)],
                             )
-                            cute.prefetch(
-                                tma_atom_sfa,
-                                tAgSFA_slice[(None, future_k_block)],
-                            )
-                            cute.prefetch(
-                                tma_atom_sfb,
-                                tBgSFB_slice[(None, future_k_block)],
-                            )
+                            if cutlass.const_expr(not self.packed_k128_scales):
+                                cute.prefetch(
+                                    tma_atom_sfa,
+                                    tAgSFA_slice[(None, future_k_block)],
+                                )
+                                cute.prefetch(
+                                    tma_atom_sfb,
+                                    tBgSFB_slice[(None, future_k_block)],
+                                )
 
                     # Peek (try_wait) AB buffer empty for k_block = prefetch_k_block_cnt + k_block + 1
                     ab_producer_state.advance()
@@ -1028,6 +1401,210 @@ def kernel(
             #
             ab_pipeline.producer_tail(ab_producer_state)
 
+        # Load and transpose packed scales on a dedicated warp.
+        if cutlass.const_expr(self.packed_k128_scales):
+            if warp_idx == self.scale_warp_id:
+                tile_sched = utils.StaticPersistentTileScheduler.create(
+                    tile_sched_params, cute.arch.block_idx(),
+                    cute.arch.grid_dim())
+                work_tile = tile_sched.initial_work_tile_info()
+                scale_stage_state = pipeline.make_pipeline_state(
+                    pipeline.PipelineUserType.Producer, self.num_scale_stage)
+                if cutlass.const_expr(self.tma_packed_scales):
+                    scale_tma_producer_state = pipeline.make_pipeline_state(
+                        pipeline.PipelineUserType.Producer,
+                        self.num_scale_stage)
+                    scale_tma_consumer_state = pipeline.make_pipeline_state(
+                        pipeline.PipelineUserType.Consumer,
+                        self.num_scale_stage)
+                scale_group_cnt = cute.ceil_div(k_block_cnt, _SCALES_PER_WORD)
+
+                while work_tile.is_valid_tile:
+                    cur_tile_coord = self._swizzle_work_tile(
+                        work_tile.tile_idx, mC_mnl)
+                    mma_tile_coord_mnl = (
+                        cur_tile_coord[0] // cute.size(tiled_mma.thr_id.shape),
+                        cur_tile_coord[1],
+                        cur_tile_coord[2],
+                    )
+                    valid_mma_n = cutlass.Int32(self.mma_tiler[1])
+                    if cutlass.const_expr(self.dynamic_mma_n):
+                        remaining_n = (
+                            cute.size(mB_nkl, mode=[0]) -
+                            mma_tile_coord_mnl[1] * self.mma_tiler[1])
+                        if remaining_n < valid_mma_n:
+                            valid_mma_n = cutlass.Int32(remaining_n)
+
+                    scale_stage_state.reset_count()
+                    if cutlass.const_expr(self.tma_packed_scales):
+                        scale_tma_producer_state.reset_count()
+                        scale_tma_consumer_state.reset_count()
+
+                    for scale_group in cutlass.range(0,
+                                                     scale_group_cnt,
+                                                     1,
+                                                     unroll=1):
+                        if cutlass.const_expr(not self.tma_packed_scales):
+                            scale_pipeline.producer_acquire(scale_stage_state)
+                        lane_idx = tidx % 32
+                        scale_stage = scale_stage_state.index
+                        if cutlass.const_expr(self.tma_packed_scales):
+                            scale_tma_pipeline.producer_acquire(
+                                scale_tma_producer_state)
+                            scale_tma_bar = (
+                                scale_tma_pipeline.producer_get_barrier(
+                                    scale_tma_producer_state))
+                            cute.copy(
+                                tma_atom_sfa,
+                                tAgSFA_tma[(None, cur_tile_coord[0],
+                                            scale_group,
+                                            mma_tile_coord_mnl[2])],
+                                tAsSFA_tma[(None,
+                                            scale_tma_producer_state.index)],
+                                tma_bar_ptr=scale_tma_bar,
+                            )
+                            cute.copy(
+                                tma_atom_sfb,
+                                tBgSFB_tma[(None, mma_tile_coord_mnl[1],
+                                            scale_group,
+                                            mma_tile_coord_mnl[2])],
+                                tBsSFB_tma[(None,
+                                            scale_tma_producer_state.index)],
+                                tma_bar_ptr=scale_tma_bar,
+                            )
+                            scale_tma_producer_state.advance()
+                            scale_tma_pipeline.consumer_wait(
+                                scale_tma_consumer_state)
+                            scale_stage = scale_tma_consumer_state.index
+                            packed_sfa = cute.make_fragment((4, ),
+                                                            cutlass.Uint32)
+                            for row_group_iter in cutlass.range_constexpr(4):
+                                tma_scale_row = lane_idx + row_group_iter * 32
+                                packed_sfa[row_group_iter] = cutlass.Uint32(
+                                    sSFA_raw[(tma_scale_row, scale_stage)])
+                            cute.arch.sync_warp()
+                            sfa_store_addr = (
+                                sSFA_raw.iterator + lane_idx * 4 + scale_stage *
+                                self.cta_tile_shape_mnk[0]).toint()
+                            sfa_store_ptr = cute.make_ptr(
+                                cutlass.Uint32,
+                                sfa_store_addr,
+                                cute.AddressSpace.smem,
+                                assumed_align=16,
+                            )
+                            sfa_store = cute.make_tensor(
+                                sfa_store_ptr, cute.make_layout((4, )))
+                            cute.autovec_copy(packed_sfa, sfa_store)
+
+                            for row_block in cutlass.range(cute.ceil_div(
+                                    self.cta_tile_shape_mnk[1], 128),
+                                                           unroll_full=True):
+                                packed_sfb = cute.make_fragment((4, ),
+                                                                cutlass.Uint32)
+                                packed_sfb.fill(0)
+                                for row_group_iter in cutlass.range_constexpr(
+                                        4):
+                                    tma_scale_row = (row_block * 128 +
+                                                     lane_idx +
+                                                     row_group_iter * 32)
+                                    if tma_scale_row < self.cta_tile_shape_mnk[
+                                            1]:
+                                        packed_sfb[
+                                            row_group_iter] = cutlass.Uint32(
+                                                sSFB_raw[(tma_scale_row,
+                                                          scale_stage)])
+                                cute.arch.sync_warp()
+                                sfb_store_addr = (
+                                    sSFB_raw.iterator + row_block * 128 +
+                                    lane_idx * 4 +
+                                    scale_stage * num_sfb_rows).toint()
+                                sfb_store_ptr = cute.make_ptr(
+                                    cutlass.Uint32,
+                                    sfb_store_addr,
+                                    cute.AddressSpace.smem,
+                                    assumed_align=16,
+                                )
+                                sfb_store = cute.make_tensor(
+                                    sfb_store_ptr, cute.make_layout((4, )))
+                                cute.autovec_copy(packed_sfb, sfb_store)
+                        else:
+                            packed_sfa = cute.make_fragment((4, ),
+                                                            cutlass.Uint32)
+                            packed_sfa.fill(0)
+                            for row_group in cutlass.range_constexpr(4):
+                                direct_scale_row = lane_idx + row_group * 32
+                                global_scale_row = (cur_tile_coord[0] *
+                                                    self.cta_tile_shape_mnk[0] +
+                                                    direct_scale_row)
+                                if global_scale_row < mSFA_mkl.shape[0]:
+                                    packed_sfa[row_group] = cutlass.Uint32(
+                                        mSFA_mkl[(
+                                            global_scale_row,
+                                            scale_group,
+                                            mma_tile_coord_mnl[2],
+                                        )])
+                            sfa_store_addr = (
+                                sSFA_raw.iterator + lane_idx * 4 + scale_stage *
+                                self.cta_tile_shape_mnk[0]).toint()
+                            sfa_store_ptr = cute.make_ptr(
+                                cutlass.Uint32,
+                                sfa_store_addr,
+                                cute.AddressSpace.smem,
+                                assumed_align=16,
+                            )
+                            sfa_store = cute.make_tensor(
+                                sfa_store_ptr, cute.make_layout((4, )))
+                            cute.autovec_copy(packed_sfa, sfa_store)
+
+                            for row_block in cutlass.range(cute.ceil_div(
+                                    self.cta_tile_shape_mnk[1], 128),
+                                                           unroll_full=True):
+                                packed_sfb = cute.make_fragment((4, ),
+                                                                cutlass.Uint32)
+                                packed_sfb.fill(0)
+                                for row_group in cutlass.range_constexpr(4):
+                                    direct_scale_row = (row_block * 128 +
+                                                        lane_idx +
+                                                        row_group * 32)
+                                    if direct_scale_row < valid_mma_n:
+                                        global_scale_row = (
+                                            mma_tile_coord_mnl[1] *
+                                            self.mma_tiler[1] +
+                                            direct_scale_row)
+                                        packed_sfb[row_group] = cutlass.Uint32(
+                                            mSFB_nkl[(
+                                                global_scale_row,
+                                                scale_group,
+                                                mma_tile_coord_mnl[2],
+                                            )])
+                                sfb_store_addr = (
+                                    sSFB_raw.iterator + row_block * 128 +
+                                    lane_idx * 4 +
+                                    scale_stage * num_sfb_rows).toint()
+                                sfb_store_ptr = cute.make_ptr(
+                                    cutlass.Uint32,
+                                    sfb_store_addr,
+                                    cute.AddressSpace.smem,
+                                    assumed_align=16,
+                                )
+                                sfb_store = cute.make_tensor(
+                                    sfb_store_ptr, cute.make_layout((4, )))
+                                cute.autovec_copy(packed_sfb, sfb_store)
+
+                        cute.arch.fence_view_async_shared()
+                        scale_pipeline.producer_commit(scale_stage_state)
+                        if cutlass.const_expr(self.tma_packed_scales):
+                            scale_tma_consumer_state.advance()
+                        scale_stage_state.advance()
+
+                    tile_sched.advance_to_next_work()
+                    work_tile = tile_sched.get_current_work()
+
+                if cutlass.const_expr(self.tma_packed_scales):
+                    scale_tma_pipeline.producer_tail(scale_tma_producer_state)
+                if cutlass.const_expr(not self.tma_packed_scales):
+                    scale_pipeline.producer_tail(scale_stage_state)
+
         #
         # Specialized MMA warp
         #
@@ -1099,18 +1676,34 @@ def kernel(
 
             ab_consumer_state = pipeline.make_pipeline_state(
                 pipeline.PipelineUserType.Consumer, self.num_ab_stage)
+            scale_consumer_state = None
+            if cutlass.const_expr(self.packed_k128_scales):
+                scale_consumer_state = pipeline.make_pipeline_state(
+                    pipeline.PipelineUserType.Consumer, self.num_scale_stage)
+                if cutlass.const_expr(self.tma_packed_scales):
+                    scale_tma_mma_state = pipeline.make_pipeline_state(
+                        pipeline.PipelineUserType.Consumer,
+                        self.num_scale_stage)
             acc_producer_state = pipeline.make_pipeline_state(
                 pipeline.PipelineUserType.Producer, self.num_acc_stage)
 
             while work_tile.is_valid_tile:
                 # Get tile coord from tile scheduler
-                cur_tile_coord = work_tile.tile_idx
+                cur_tile_coord = self._swizzle_work_tile(
+                    work_tile.tile_idx, mC_mnl)
                 mma_tile_coord_mnl = (
                     cur_tile_coord[0] // cute.size(tiled_mma.thr_id.shape),
                     cur_tile_coord[1],
                     cur_tile_coord[2],
                 )
 
+                valid_mma_n = cutlass.Int32(self.mma_tiler[1])
+                if cutlass.const_expr(self.dynamic_mma_n):
+                    remaining_n = (cute.size(mB_nkl, mode=[0]) -
+                                   mma_tile_coord_mnl[1] * self.mma_tiler[1])
+                    if remaining_n < valid_mma_n:
+                        valid_mma_n = cutlass.Int32(remaining_n)
+
                 if cutlass.const_expr(self.overlapping_accum):
                     acc_stage_index = acc_producer_state.phase ^ 1
                 else:
@@ -1126,7 +1719,18 @@ def kernel(
                 if ab_consumer_state.count < k_block_cnt and is_leader_cta:
                     peek_ab_full_status = ab_pipeline.consumer_try_wait(
                         ab_consumer_state)
-
+                if cutlass.const_expr(self.packed_k128_scales):
+                    scale_group_cnt = cute.ceil_div(k_block_cnt,
+                                                    _SCALES_PER_WORD)
+                    scale_consumer_state.reset_count()
+                    if cutlass.const_expr(self.tma_packed_scales):
+                        scale_tma_mma_state.reset_count()
+                    peek_scale_full_status = cutlass.Boolean(1)
+                    if (scale_consumer_state.count < scale_group_cnt
+                            and is_leader_cta):
+                        peek_scale_full_status = (
+                            scale_pipeline.consumer_try_wait(
+                                scale_consumer_state))
                 #
                 # Wait for accumulator buffer empty
                 #
@@ -1134,7 +1738,8 @@ def kernel(
                     acc_pipeline.producer_acquire(acc_producer_state)
 
                 tCtSFB_mma = tCtSFB
-                if cutlass.const_expr(self.cta_tile_shape_mnk[1] == 192):
+                if cutlass.const_expr(not self.packed_k128_scales
+                                      and self.cta_tile_shape_mnk[1] == 192):
                     # If this is an ODD tile, shift the TMEM start address for cta_tile_shape_n=192 case by two words (ignores first 64 columns of SFB)
                     offset = cutlass.Int32(
                         2) if mma_tile_coord_mnl[1] % 2 == 1 else cutlass.Int32(
@@ -1145,7 +1750,8 @@ def kernel(
                         dtype=self.sf_dtype,
                     )
                     tCtSFB_mma = cute.make_tensor(shifted_ptr, tCtSFB_layout)
-                elif cutlass.const_expr(self.cta_tile_shape_mnk[1] == 64):
+                elif cutlass.const_expr(not self.packed_k128_scales
+                                        and self.cta_tile_shape_mnk[1] == 64):
                     # Move in increments of 64 columns of SFB
                     offset = cutlass.Int32((mma_tile_coord_mnl[1] % 2) * 2)
                     shifted_ptr = cute.recast_ptr(
@@ -1163,71 +1769,140 @@ def kernel(
                 #
                 # Mma mainloop
                 #
-                for k_block in range(k_block_cnt):
+                for k_block in cutlass.range(0, k_block_cnt, 1, unroll=4):
                     if is_leader_cta:
-                        # Conditionally wait for AB buffer full
                         ab_pipeline.consumer_wait(ab_consumer_state,
                                                   peek_ab_full_status)
+                        if cutlass.const_expr(self.packed_k128_scales):
+                            if k_block % _SCALES_PER_WORD == 0:
+                                scale_pipeline.consumer_wait(
+                                    scale_consumer_state,
+                                    peek_scale_full_status)
 
                         #  Copy SFA/SFB from smem to tmem
+                        s2t_stage = ab_consumer_state.index
+                        if cutlass.const_expr(self.packed_k128_scales):
+                            s2t_stage = scale_consumer_state.index
                         s2t_stage_coord = (
                             None,
                             None,
                             None,
                             None,
-                            ab_consumer_state.index,
+                            s2t_stage,
                         )
                         tCsSFA_compact_s2t_staged = tCsSFA_compact_s2t[
                             s2t_stage_coord]
                         tCsSFB_compact_s2t_staged = tCsSFB_compact_s2t[
                             s2t_stage_coord]
-                        cute.copy(
-                            tiled_copy_s2t_sfa,
-                            tCsSFA_compact_s2t_staged,
-                            tCtSFA_compact_s2t,
-                        )
-                        cute.copy(
-                            tiled_copy_s2t_sfb,
-                            tCsSFB_compact_s2t_staged,
-                            tCtSFB_compact_s2t,
-                        )
-
-                        # tCtAcc += tCrA * tCrSFA * tCrB * tCrSFB
-                        num_kphases = cute.size(tCrA, mode=[2])
-                        for kphase_idx in cutlass.range(num_kphases,
-                                                        unroll_full=True):
-                            kphase_coord = (
-                                None,
-                                None,
-                                kphase_idx,
-                                ab_consumer_state.index,
+                        if cutlass.const_expr(self.packed_k128_scales):
+                            if k_block % _SCALES_PER_WORD == 0:
+                                cute.copy(
+                                    tiled_copy_s2t_sfa,
+                                    tCsSFA_compact_s2t_staged,
+                                    tCtSFA_compact_s2t,
+                                )
+                                cute.copy(
+                                    tiled_copy_s2t_sfb,
+                                    tCsSFB_compact_s2t_staged,
+                                    tCtSFB_compact_s2t,
+                                )
+                        else:
+                            cute.copy(
+                                tiled_copy_s2t_sfa,
+                                tCsSFA_compact_s2t_staged,
+                                tCtSFA_compact_s2t,
                             )
-
-                            # Set SFA/SFB tensor to tiled_mma
-                            sf_kphase_coord = (None, None, kphase_idx)
-                            tiled_mma.set(
-                                tcgen05.Field.SFA,
-                                tCtSFA[sf_kphase_coord].iterator,
-                            )
-                            tiled_mma.set(
-                                tcgen05.Field.SFB,
-                                tCtSFB_mma[sf_kphase_coord].iterator,
+                            cute.copy(
+                                tiled_copy_s2t_sfb,
+                                tCsSFB_compact_s2t_staged,
+                                tCtSFB_compact_s2t,
                             )
 
-                            cute.gemm(
-                                tiled_mma,
-                                tCtAcc,
-                                tCrA[kphase_coord],
-                                tCrB[kphase_coord],
-                                tCtAcc,
+                        if cutlass.const_expr(self.dynamic_mma_n):
+                            # Mask the swapped N-tail in the MMA descriptor.
+                            dynamic_mxf8_mainloop.issue_dynamic_mxf8_mma_tile(
+                                acc_tensor=tCtAcc,
+                                a_frag_tile=tCrA[(None, None, None,
+                                                  ab_consumer_state.index)],
+                                b_frag_tile=tCrB[(None, None, None,
+                                                  ab_consumer_state.index)],
+                                sfa_tensor=tCtSFA,
+                                sfb_tensor=tCtSFB_mma,
+                                k_tile_idx=k_block,
+                                valid_n=valid_mma_n,
+                                sfa_kphase_idx=(k_block % _SCALES_PER_WORD
+                                                if self.packed_k128_scales else
+                                                None),
+                                sfb_kphase_idx=(k_block % _SCALES_PER_WORD
+                                                if self.packed_k128_scales else
+                                                None),
+                                mma_tiler_mnk=self.mma_tiler,
                             )
-
-                            # Enable accumulate on tCtAcc after first kphase
-                            tiled_mma.set(tcgen05.Field.ACCUMULATE, True)
+                        else:
+                            num_kphases = cute.size(tCrA, mode=[2])
+                            for kphase_idx in cutlass.range(num_kphases,
+                                                            unroll_full=True):
+                                kphase_coord = (
+                                    None,
+                                    None,
+                                    kphase_idx,
+                                    ab_consumer_state.index,
+                                )
+
+                                sf_kphase_idx = kphase_idx
+                                if cutlass.const_expr(self.packed_k128_scales):
+                                    sf_kphase_idx = k_block % _SCALES_PER_WORD
+                                sf_kphase_coord = (None, None, sf_kphase_idx)
+                                tiled_mma.set(
+                                    tcgen05.Field.SFA,
+                                    tCtSFA[sf_kphase_coord].iterator,
+                                )
+                                tiled_mma.set(
+                                    tcgen05.Field.SFB,
+                                    tCtSFB_mma[sf_kphase_coord].iterator,
+                                )
+
+                                cute.gemm(
+                                    tiled_mma,
+                                    tCtAcc,
+                                    tCrA[kphase_coord],
+                                    tCrB[kphase_coord],
+                                    tCtAcc,
+                                )
+
+                                tiled_mma.set(tcgen05.Field.ACCUMULATE, True)
 
                         # Async arrive AB buffer empty
                         ab_pipeline.consumer_release(ab_consumer_state)
 
+                        if cutlass.const_expr(self.packed_k128_scales):
+                            if k_block % _SCALES_PER_WORD == 0:
+                                if cutlass.const_expr(self.tma_packed_scales):
+                                    scale_empty_mbar = (
+                                        scale_tma_pipeline.sync_object_empty.
+                                        get_barrier(scale_tma_mma_state.index))
+                                else:
+                                    scale_empty_mbar = (
+                                        scale_pipeline.sync_object_empty.
+                                        get_barrier(scale_consumer_state.index))
+                                with cute.arch.elect_one():
+                                    if cutlass.const_expr(use_2cta_instrs):
+                                        tcgen05.commit(
+                                            scale_empty_mbar,
+                                            ab_pipeline.consumer_mask,
+                                            self.cta_group,
+                                        )
+                                    else:
+                                        tcgen05.commit(scale_empty_mbar)
+                                scale_consumer_state.advance()
+                                if cutlass.const_expr(self.tma_packed_scales):
+                                    scale_tma_mma_state.advance()
+                                peek_scale_full_status = cutlass.Boolean(1)
+                                if scale_consumer_state.count < scale_group_cnt:
+                                    peek_scale_full_status = (
+                                        scale_pipeline.consumer_try_wait(
+                                            scale_consumer_state))
+
                     # Peek (try_wait) AB buffer full for k_block = k_block + 1
                     ab_consumer_state.advance()
                     peek_ab_full_status = cutlass.Boolean(1)
@@ -1327,7 +2002,8 @@ def kernel(
             while work_tile.is_valid_tile:
 
                 # Get tile coord from tile scheduler
-                cur_tile_coord = work_tile.tile_idx
+                cur_tile_coord = self._swizzle_work_tile(
+                    work_tile.tile_idx, mC_mnl)
                 mma_tile_coord_mnl = (
                     cur_tile_coord[0] // cute.size(tiled_mma.thr_id.shape),
                     cur_tile_coord[1],
@@ -1399,8 +2075,10 @@ def kernel(
                     # Convert to C type
                     #
                     acc_vec = tiled_copy_r2s.retile(tTR_rAcc).load()
-                    acc_vec = epilogue_op(alpha_value *
-                                          acc_vec.to(self.c_dtype))
+                    acc_vec = acc_vec.to(self.c_dtype)
+                    if cutlass.const_expr(self.apply_alpha):
+                        acc_vec = alpha_value * acc_vec
+                    acc_vec = epilogue_op(acc_vec)
                     tRS_rC.store(acc_vec)
 
                     #
diff --git a/tensorrt_llm/_torch/cute_dsl_kernels/blackwell/dynamic_mxf8_mainloop.py b/tensorrt_llm/_torch/cute_dsl_kernels/blackwell/dynamic_mxf8_mainloop.py
new file mode 100644
index 000000000000..974e7074e1cd
--- /dev/null
+++ b/tensorrt_llm/_torch/cute_dsl_kernels/blackwell/dynamic_mxf8_mainloop.py
@@ -0,0 +1,133 @@
+# SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: Apache-2.0
+"""Runtime-N tcgen05 block-scaled MXF8 MMA helpers for SM100."""
+
+from typing import Optional
+
+import cutlass.cute as cute
+from cutlass._mlir import ir
+from cutlass._mlir.dialects import builtin, llvm
+from cutlass._mlir.dialects import nvvm as _nvvm_raw
+from cutlass.cute.arch.nvvm_wrappers import nvvm
+from cutlass.cutlass_dsl import Boolean, Int32, Uint32, dsl_user_op
+
+_MXF8_SCALE_VEC_SIZE = _nvvm_raw.Tcgen05MMAScaleVecSize.X1
+
+
+def _align16(x):
+    return (Int32(x) + Int32(15)) & Int32(-16)
+
+
+@dsl_user_op
+def compute_non_leader_cta_load_shift(
+    *,
+    valid_n,
+    mma_tiler_n: int,
+    loc: Optional[ir.Location] = None,
+    ip: Optional[ir.InsertionPoint] = None,
+) -> Int32:
+    return (_align16(valid_n) >> Int32(1)) - Int32(mma_tiler_n // 2)
+
+
+def _build_static_idesc_base(umma_m: int) -> int:
+    # MXF8 E4M3 operands use format 0, UE8M0 scales use scale format 1.
+    return (((umma_m >> 4) & 0x1F) << 24) | (1 << 23)
+
+
+def _as_value(iterator) -> ir.Value:
+    return iterator.value if hasattr(iterator, "value") else iterator
+
+
+def _smem_desc_to_i64(value: ir.Value) -> ir.Value:
+    return builtin.unrealized_conversion_cast([ir.IntegerType.get_signless(64)], [value])
+
+
+def _tmem_ptr_to_i32(value: ir.Value) -> ir.Value:
+    return builtin.unrealized_conversion_cast([ir.IntegerType.get_signless(32)], [value])
+
+
+def _i32_to_tmem_ptr(value: ir.Value) -> ir.Value:
+    ptr_ty = llvm.PointerType.get(cute.AddressSpace.tmem.value)
+    return llvm.inttoptr(ptr_ty, value)
+
+
+@dsl_user_op
+def issue_dynamic_mxf8_mma_tile(
+    *,
+    acc_tensor,
+    a_frag_tile,
+    b_frag_tile,
+    sfa_tensor,
+    sfb_tensor,
+    k_tile_idx,
+    valid_n,
+    sfa_kphase_idx=None,
+    sfb_kphase_idx=None,
+    mma_tiler_mnk: tuple = (256, 208, 128),
+    loc: Optional[ir.Location] = None,
+    ip: Optional[ir.InsertionPoint] = None,
+) -> None:
+    """Issue one K=128 MXF8 tile with a runtime UMMA N extent."""
+    static_base = _build_static_idesc_base(mma_tiler_mnk[0])
+    n_dim = _align16(valid_n) >> Int32(3)
+    num_k_inner = mma_tiler_mnk[2] // 32
+    group_kind = (
+        _nvvm_raw.Tcgen05GroupKind
+        if hasattr(_nvvm_raw, "Tcgen05GroupKind")
+        else _nvvm_raw.CTAGroupKind
+    )
+
+    acc_atom = acc_tensor[(None, 0, 0)]
+    acc_i32 = _tmem_ptr_to_i32(_as_value(acc_atom.iterator))
+    idesc_base = Uint32(static_base) | (Uint32(n_dim) << Uint32(17))
+
+    packed_sfa_i32 = None
+    packed_sfb_i32 = None
+    if sfa_kphase_idx is not None:
+        packed_sfa_atom = sfa_tensor[(None, 0, sfa_kphase_idx)]
+        packed_sfa_i32 = _tmem_ptr_to_i32(_as_value(packed_sfa_atom.iterator))
+    if sfb_kphase_idx is not None:
+        packed_sfb_atom = sfb_tensor[(None, 0, sfb_kphase_idx)]
+        packed_sfb_i32 = _tmem_ptr_to_i32(_as_value(packed_sfb_atom.iterator))
+
+    packed_idesc = idesc_base
+    # Encode packed-scale TMEM phases in the instruction descriptor.
+    if packed_sfa_i32 is not None:
+        packed_idesc = packed_idesc | ((Uint32(packed_sfa_i32) & Uint32(0xC0000000)) >> Uint32(1))
+    if packed_sfb_i32 is not None:
+        packed_idesc = packed_idesc | ((Uint32(packed_sfb_i32) & Uint32(0xC0000000)) >> Uint32(26))
+
+    for k_inner in range(num_k_inner):
+        a_atom = a_frag_tile[(None, 0, k_inner)]
+        b_atom = b_frag_tile[(None, 0, k_inner)]
+        sfa_idx = k_inner if sfa_kphase_idx is None else sfa_kphase_idx
+        sfb_idx = k_inner if sfb_kphase_idx is None else sfb_kphase_idx
+
+        operand_a = _smem_desc_to_i64(_as_value(a_atom.iterator))
+        operand_b = _smem_desc_to_i64(_as_value(b_atom.iterator))
+        sfa_i32 = packed_sfa_i32
+        sfb_i32 = packed_sfb_i32
+        idesc = packed_idesc
+        if sfa_i32 is None:
+            sfa_atom = sfa_tensor[(None, 0, sfa_idx)]
+            sfa_i32 = _tmem_ptr_to_i32(_as_value(sfa_atom.iterator))
+            idesc = idesc | ((Uint32(sfa_i32) & Uint32(0xC0000000)) >> Uint32(1))
+        if sfb_i32 is None:
+            sfb_atom = sfb_tensor[(None, 0, sfb_idx)]
+            sfb_i32 = _tmem_ptr_to_i32(_as_value(sfb_atom.iterator))
+            idesc = idesc | ((Uint32(sfb_i32) & Uint32(0xC0000000)) >> Uint32(26))
+        accum = k_tile_idx != 0 if k_inner == 0 else True
+
+        with cute.arch.elect_one():
+            nvvm.tcgen05_mma_block_scale(
+                mma_kind=_nvvm_raw.Tcgen05MMAKind.MXF8F6F4,
+                cta_group=(group_kind.CTA_2 if mma_tiler_mnk[0] == 256 else group_kind.CTA_1),
+                d=_i32_to_tmem_ptr(acc_i32),
+                a=operand_a,
+                b=operand_b,
+                idesc=idesc.ir_value(),
+                enable_input_d=Boolean(accum).ir_value(),
+                scale_a=_i32_to_tmem_ptr(sfa_i32),
+                scale_b=_i32_to_tmem_ptr(sfb_i32),
+                scale_vec_size=_MXF8_SCALE_VEC_SIZE,
+            )
diff --git a/tensorrt_llm/_torch/modules/attention.py b/tensorrt_llm/_torch/modules/attention.py
index 7b7a19e046ac..ba512ba10c7a 100644
--- a/tensorrt_llm/_torch/modules/attention.py
+++ b/tensorrt_llm/_torch/modules/attention.py
@@ -24,6 +24,7 @@
 from ..attention_backend.sparse.dsa import (
     DSAtrtllmAttentionMetadata, transform_local_topk_and_prepare_pool_view)
 from ..attention_backend.utils import create_attention, get_attention_backend
+from ..cute_dsl_utils import IS_CUTLASS_DSL_AVAILABLE
 from ..distributed import (AllReduceParams, HelixAllToAllNative, alltoall_helix,
                            cp_allgather, reducescatter)
 from ..model_config import ModelConfig
@@ -47,6 +48,29 @@ def _is_env_truthy(name: str) -> bool:
     return os.environ.get(name, "").strip().lower() in ("1", "true", "on")
 
 
+_DSV4_OB_SCALE_BLOCK_K = 128
+_UE8M0_SCALES_PER_WORD = 4
+_DSV4_OB_PACKED_SCALE_K = _DSV4_OB_SCALE_BLOCK_K * _UE8M0_SCALES_PER_WORD
+_DSV4_OB_SPLIT_K_MAX_M = 128
+_DSV4_OB_CUTE_SK1_MIN_M = 256
+_DSV4_PRO_OB_N = 7168
+_DSV4_PRO_OB_K = 16384
+
+
+def _select_dsv4_ob_split_k(num_tokens: int,
+                            configured_split: Optional[int] = None) -> int:
+    if configured_split is None:
+        env_split = int(os.environ.get("TRTLLM_DSV4_OB_SPLIT_K", "0"))
+        split_k = env_split if env_split else (
+            2 if 0 < num_tokens <= _DSV4_OB_SPLIT_K_MAX_M else 1)
+    else:
+        split_k = configured_split
+    if split_k not in (1, 2, 4):
+        raise ValueError(
+            f"unsupported DeepSeek-V4 O_b split-K factor: {split_k}")
+    return split_k
+
+
 def extract_extra_attrs(layer_idx: str, attn_type: str):
     assert attn_type in ["mla", "attn"], "Invalid attention type"
     extra_attrs = get_model_extra_attrs()
@@ -2008,6 +2032,117 @@ def _validate_dsv4_epilogue_buffers(
                 "Invalid DSv4 fused epilogue buffers for current token count.")
         return fp8_o, output_sf
 
+    def _should_use_fused_oproj(self) -> bool:
+        # Eligible Blackwell paths are on by default; keep an emergency fallback.
+        return (not _is_env_truthy("TRTLLM_DSV4_DISABLE_FUSED_OPROJ")
+                and is_sm_100f() and IS_CUTLASS_DSL_AVAILABLE
+                and self.n_local_groups == self.num_groups
+                and getattr(self.o_b_proj, "tp_size", 1) == 1
+                and self.o_a_proj.dtype == torch.float8_e4m3fn
+                and self.o_b_proj.has_fp8_block_scales and not getattr(
+                    self.o_b_proj, "use_cute_dsl_blockscaling_mm", False))
+
+    def _fused_oa_ob_proj(self, attn_fp8: torch.Tensor,
+                          attn_scale: torch.Tensor,
+                          num_tokens: int) -> torch.Tensor:
+        num_groups, o_lora_rank = self.n_local_groups, self.o_lora_rank
+        aligned_m = ((num_tokens + _UE8M0_SCALES_PER_WORD - 1) //
+                     _UE8M0_SCALES_PER_WORD) * _UE8M0_SCALES_PER_WORD
+        n_tiles = ((o_lora_rank + _DSV4_OB_SCALE_BLOCK_K - 1) //
+                   _DSV4_OB_SCALE_BLOCK_K)
+        packed_k = ((num_groups * n_tiles + _UE8M0_SCALES_PER_WORD - 1) //
+                    _UE8M0_SCALES_PER_WORD)
+        o_lora_fp8 = torch.empty([num_tokens, num_groups, o_lora_rank],
+                                 device=attn_fp8.device,
+                                 dtype=torch.float8_e4m3fn)
+        sf_storage = torch.zeros(packed_k * aligned_m,
+                                 device=attn_fp8.device,
+                                 dtype=torch.int32)
+        sf_out = torch.as_strided(sf_storage, (num_tokens, packed_k),
+                                  (1, aligned_m))
+        torch.ops.trtllm.cute_dsl_fp8_bmm_blackwell_fp8out(
+            attn_fp8, self.o_a_proj, attn_scale, self.o_a_proj_scale,
+            o_lora_fp8.transpose(0, 1), sf_out)
+        return self._fused_ob_gemm(o_lora_fp8.flatten(1), sf_out, num_tokens)
+
+    def _fused_ob_gemm(self, o_lora_fp8: torch.Tensor, sf_out: torch.Tensor,
+                       num_tokens: int) -> torch.Tensor:
+        """Run the DSV4-Pro O_b FP8 GEMM and emit mHC-ready split partials."""
+        from tensorrt_llm import deep_gemm
+
+        M, N = num_tokens, self.hidden_size
+        K_ob = o_lora_fp8.shape[1]
+        w, wsf = self.o_b_proj.weight, self.o_b_proj.weight_scale
+        SK = _select_dsv4_ob_split_k(M, getattr(self, "ob_split_k", None))
+        packed_k = K_ob // _DSV4_OB_PACKED_SCALE_K
+        # CuTe specializes the DSV4-Pro O_b shape.
+        use_cute_tactic = (IS_CUTLASS_DSL_AVAILABLE and M > 0
+                           and (M >= _DSV4_OB_CUTE_SK1_MIN_M or SK > 1)
+                           and N == _DSV4_PRO_OB_N and K_ob == _DSV4_PRO_OB_K
+                           and K_ob % (_DSV4_OB_PACKED_SCALE_K * SK) == 0)
+        activation_layout_ok = (sf_out.dtype == torch.int32
+                                and sf_out.dim() == 2
+                                and sf_out.shape == (M, packed_k)
+                                and sf_out.stride(0) == 1
+                                and sf_out.stride(1) >= M and
+                                sf_out.stride(1) % _UE8M0_SCALES_PER_WORD == 0)
+        if use_cute_tactic and not activation_layout_ok:
+            raise RuntimeError(
+                "DSV4-Pro O_b CuTe tactic requires packed activation scales.")
+
+        cute_wsf = wsf
+        if use_cute_tactic and cute_wsf.dtype != torch.int32:
+            from tensorrt_llm.quantization.utils.fp8_utils import \
+                transform_sf_into_required_layout
+            cute_wsf = getattr(self.o_b_proj, "_ob_wsf_int", None)
+            if cute_wsf is None:
+                # Convert the static weight scale once.
+                cute_wsf = transform_sf_into_required_layout(
+                    wsf,
+                    mn=N,
+                    k=K_ob,
+                    recipe=(
+                        1,
+                        _DSV4_OB_SCALE_BLOCK_K,
+                        _DSV4_OB_SCALE_BLOCK_K,
+                    ),
+                    is_sfa=False)
+                self.o_b_proj._ob_wsf_int = cute_wsf
+
+        weight_layout_ok = (cute_wsf.dtype == torch.int32
+                            and cute_wsf.dim() == 2
+                            and cute_wsf.shape == (N, packed_k)
+                            and cute_wsf.stride(0) == 1
+                            and cute_wsf.stride(1) == N)
+        if use_cute_tactic and not weight_layout_ok:
+            raise RuntimeError(
+                "DSV4-Pro O_b CuTe tactic requires packed weight scales.")
+
+        if not use_cute_tactic:
+            # Keep DeepGEMM for shapes outside the tuned CuTe buckets.
+            hidden = torch.empty([M, N],
+                                 device=o_lora_fp8.device,
+                                 dtype=self.dtype)
+            deep_gemm.fp8_gemm_nt((o_lora_fp8, sf_out), (w, wsf),
+                                  hidden,
+                                  c=None,
+                                  disable_ue8m0_cast=False)
+            return hidden
+
+        logger.info_once(
+            f"[obsk-fused] O_b CuTe DSL ENGAGED: SK={SK} M={M} "
+            f"K={K_ob} N={N}",
+            key="obsk_fused_engaged")
+        # mHC consumes split-major rows directly.
+        partials = torch.empty([SK, M, N],
+                               device=o_lora_fp8.device,
+                               dtype=self.dtype)
+        torch.ops.trtllm.dsv4_fp8_splitk_gemm(o_lora_fp8, sf_out, w, cute_wsf,
+                                              partials, SK)
+        if SK == 1:
+            return partials[0]
+        return partials.reshape(SK * M, N)
+
     def _deepseek_v4_o_proj(
         self,
         attn_out_latent: Union[torch.Tensor, Tuple[torch.Tensor, torch.Tensor]],
@@ -2017,6 +2152,9 @@ def _deepseek_v4_o_proj(
             attn_fp8, attn_scale = attn_out_latent
             num_tokens = attn_fp8.shape[1]
 
+            if self._should_use_fused_oproj():
+                return self._fused_oa_ob_proj(attn_fp8, attn_scale, num_tokens)
+
             o_lora = torch.empty(
                 [num_tokens, self.n_local_groups, self.o_lora_rank],
                 device=attn_fp8.device,
@@ -2031,14 +2169,7 @@ def _deepseek_v4_o_proj(
         attn_out_latent = attn_out_latent.view(num_tokens, self.num_heads_tp,
                                                -1)
 
-        # When o_a_proj is FP8 on SM100 (which is always the case for DSv4
-        # under FP8 block-scales after init), fuse the inverse-RoPE into the
-        # FP8-quant epilogue (vLLM-ported Triton kernel) and call
-        # cute_dsl_fp8_bmm_blackwell directly. Saves one BF16 read+write of
-        # the latent vs the mla_rope_inplace +
-        # fp8_batched_quantize_1x128_permute102 pair. Decoupled from
-        # use_cute_dsl_blockscaling_bmm (which gates the separate K/V
-        # absorption BMM kernel choice).
+        # Fuse inverse RoPE and FP8 quantization on SM100.
         fused_inv_rope_fp8 = (self.o_a_proj.dtype == torch.float8_e4m3fn
                               and is_sm_100f())
         if fused_inv_rope_fp8:
@@ -2055,6 +2186,9 @@ def _deepseek_v4_o_proj(
                     128,
                     self.inverse_rotary_emb.is_neox,
                 ))
+            if self._should_use_fused_oproj():
+                return self._fused_oa_ob_proj(attn_fp8, attn_scale, num_tokens)
+
             o_lora = torch.empty(
                 [num_tokens, self.n_local_groups, self.o_lora_rank],
                 device=attn_out_latent.device,
diff --git a/tensorrt_llm/_torch/modules/mhc/hyper_connection.py b/tensorrt_llm/_torch/modules/mhc/hyper_connection.py
index 5c02f2f97ea4..47a0dfcde13d 100644
--- a/tensorrt_llm/_torch/modules/mhc/hyper_connection.py
+++ b/tensorrt_llm/_torch/modules/mhc/hyper_connection.py
@@ -163,7 +163,7 @@ def fused_hc(
         kernel launch on the model's pre-attention norm.
 
         Args:
-            x_prev:        [..., hidden]    bf16  (attn / MoE output of prev block)
+            x_prev:        [..., hidden] or split-major [SK * B, hidden] bf16
             residual_prev: [..., mult, hidden] bf16
             post_mix_prev: [..., mult] or [..., mult, 1] fp32
             comb_mix_prev: [..., mult, mult] fp32
@@ -190,7 +190,14 @@ def fused_hc(
 
         residual_prev_flat = residual_prev.reshape(-1, n, hidden).contiguous()
         B = residual_prev_flat.shape[0]
-        x_prev_flat = x_prev.reshape(B, hidden).contiguous()
+        expected_per_split = B * hidden
+        assert x_prev.numel() % expected_per_split == 0, (
+            f"x_prev has {x_prev.numel()} elements, expected an integer multiple "
+            f"of B * hidden = {expected_per_split}"
+        )
+        x_num_splits = x_prev.numel() // expected_per_split
+        assert x_num_splits in (1, 2, 4), f"unsupported x_prev split count: {x_num_splits}"
+        x_prev_flat = x_prev.reshape(x_num_splits * B, hidden).contiguous()
         post_mix_prev_flat = post_mix_prev.reshape(B, n)
         comb_mix_prev_flat = comb_mix_prev.reshape(B, n, n)
 
diff --git a/tensorrt_llm/_torch/modules/mhc/mhc_cuda.py b/tensorrt_llm/_torch/modules/mhc/mhc_cuda.py
index 17edc3122fdf..be3c15e389fb 100644
--- a/tensorrt_llm/_torch/modules/mhc/mhc_cuda.py
+++ b/tensorrt_llm/_torch/modules/mhc/mhc_cuda.py
@@ -660,6 +660,7 @@ def _fused_hc_call(
     sinkhorn_repeat: int,
     norm_weight: torch.Tensor | None = None,
     norm_eps: float = 0.0,
+    x_num_splits: int = 1,
 ):
     torch.ops.trtllm.mhc_fused_hc(
         x_prev,
@@ -691,6 +692,7 @@ def _fused_hc_call(
         tile_m,
         norm_weight,
         norm_eps,
+        x_num_splits,
     )
 
 
@@ -799,6 +801,7 @@ def _alloc_fused_hc_outputs(
 # valid everywhere.
 _FUSED_HC_FALLBACK_TACTIC_MMA = ("fused_half_mma", 0, 0, 0, 1)
 _FUSED_HC_FALLBACK_TACTIC_FMA = ("fused_half_fma", 2, 1, 256, 1)
+_FUSED_HC_SPLIT_X_FMA_TACTIC = ("fused_half_fma", 3, 2, 256, 1)
 
 
 def _get_fused_hc_fallback_tactic(hidden_size: int | None = None):
@@ -960,6 +963,17 @@ def forward(self, inputs, *, tactic=-1, **kwargs):
             norm_weight = norm_weight.contiguous()
 
         B = residual_prev.shape[0]
+        if x_prev.ndim != 2 or x_prev.shape[1] != self.hidden_size:
+            raise ValueError(
+                f"x_prev must have shape [SK * B, {self.hidden_size}], got {tuple(x_prev.shape)}"
+            )
+        if B <= 0 or x_prev.shape[0] % B != 0:
+            raise ValueError(
+                f"x_prev rows ({x_prev.shape[0]}) must be an integer multiple of residual rows ({B})"
+            )
+        x_num_splits = x_prev.shape[0] // B
+        if x_num_splits not in (1, 2, 4):
+            raise ValueError(f"unsupported x_prev split count: {x_num_splits}")
         (
             residual_cur,
             post_mix_cur,
@@ -1000,6 +1014,7 @@ def forward(self, inputs, *, tactic=-1, **kwargs):
             self.sinkhorn_repeat,
             norm_weight=norm_weight,
             norm_eps=norm_eps,
+            x_num_splits=x_num_splits,
         )
         return residual_cur, post_mix_cur, comb_mix_cur, layer_input_cur
 
@@ -1063,6 +1078,11 @@ def mhc_fused_hc(
 ):
     """Fuse the previous block's post_mapping with the current block's pre_mapping.
 
+    ``x_prev`` may be either reduced ``[B, hidden]`` or split-major O-projection
+    partials ``[SK * B, hidden]``. For SK2/SK4 the half-MMA path streams and
+    sums partials in its pmap x-ring, while half-FMA sums them at its register
+    x-load. Both accumulate in FP32.
+
     The autotuner chooses between four backends:
       * "fused_half_mma" — 2-kernel tcgen05 TF32 pmap+GEMM atomic + bigfuse.
       * "fused_half_fma" — 2-kernel pmap inline + FMA GEMM + sqrsum + bigfuse.
@@ -1094,15 +1114,40 @@ def mhc_fused_hc(
         sinkhorn_repeat=sinkhorn_repeat,
     )
 
-    tuner = AutoTuner.get()
-    _, best_tactic = tuner.choose_one(
-        "trtllm::mhc_fused_hc",
-        [runner],
-        MhcFusedHcRunner.tuning_config,
-        [x_prev, residual_prev, post_mix_prev, comb_mix_prev, w_t_cur, hc_scale_cur, hc_base_cur],
-        norm_weight=norm_weight,
-        norm_eps=norm_eps,
-    )
+    B = residual_prev.shape[0]
+    if x_prev.ndim != 2 or B <= 0 or x_prev.shape[0] % B != 0:
+        raise ValueError(
+            f"x_prev must have split-major shape [SK * B, hidden], got {tuple(x_prev.shape)} "
+            f"for B={B}"
+        )
+    x_num_splits = x_prev.shape[0] // B
+    if x_num_splits not in (1, 2, 4):
+        raise ValueError(f"unsupported x_prev split count: {x_num_splits}")
+    if x_num_splits in (2, 4):
+        if B <= 32:
+            # CUDA-core FMA wins at tiny M.
+            best_tactic = _FUSED_HC_SPLIT_X_FMA_TACTIC
+        else:
+            # C++ selects split-K and BigFuse geometry from M.
+            best_tactic = _FUSED_HC_FALLBACK_TACTIC_MMA
+    else:
+        tuner = AutoTuner.get()
+        _, best_tactic = tuner.choose_one(
+            "trtllm::mhc_fused_hc",
+            [runner],
+            MhcFusedHcRunner.tuning_config,
+            [
+                x_prev,
+                residual_prev,
+                post_mix_prev,
+                comb_mix_prev,
+                w_t_cur,
+                hc_scale_cur,
+                hc_base_cur,
+            ],
+            norm_weight=norm_weight,
+            norm_eps=norm_eps,
+        )
 
     return runner(
         inputs=[
diff --git a/tests/unittest/_torch/attention/sparse/deepseek_v4/test_deepseek_v4_o_proj.py b/tests/unittest/_torch/attention/sparse/deepseek_v4/test_deepseek_v4_o_proj.py
index ee746b64e84a..15132f67ff7f 100644
--- a/tests/unittest/_torch/attention/sparse/deepseek_v4/test_deepseek_v4_o_proj.py
+++ b/tests/unittest/_torch/attention/sparse/deepseek_v4/test_deepseek_v4_o_proj.py
@@ -9,10 +9,12 @@
 from _torch.helpers import per_block_cast_to_fp8_e8m0, per_token_cast_to_fp8_e8m0
 from utils.util import skip_pre_blackwell
 
+import tensorrt_llm._torch.modules.attention as attention_module
 from tensorrt_llm._torch.attention_backend.interface import PositionalEmbeddingParams, RopeParams
+from tensorrt_llm._torch.custom_ops import cute_dsl_custom_ops
 from tensorrt_llm._torch.model_config import ModelConfig
 from tensorrt_llm._torch.models.modeling_deepseekv3 import weight_dequant
-from tensorrt_llm._torch.modules.attention import MLA
+from tensorrt_llm._torch.modules.attention import MLA, _select_dsv4_ob_split_k
 from tensorrt_llm._utils import get_sm_version
 from tensorrt_llm.functional import PositionEmbeddingType
 from tensorrt_llm.llmapi.llm_args import DeepSeekV4SparseAttentionConfig
@@ -25,6 +27,280 @@
 FP8_O_PROJ_DIFF_TOL = 2e-3
 
 
+@pytest.mark.parametrize(
+    ("num_tokens", "expected_split"),
+    [(1, 2), (16, 2), (32, 2), (64, 2), (96, 2), (128, 2), (160, 1), (16384, 1)],
+)
+def test_select_dsv4_ob_split_k_auto_policy(num_tokens, expected_split, monkeypatch):
+    monkeypatch.delenv("TRTLLM_DSV4_OB_SPLIT_K", raising=False)
+    assert _select_dsv4_ob_split_k(num_tokens) == expected_split
+
+
+def test_select_dsv4_ob_split_k_overrides(monkeypatch):
+    monkeypatch.setenv("TRTLLM_DSV4_OB_SPLIT_K", "4")
+    assert _select_dsv4_ob_split_k(32) == 4
+    assert _select_dsv4_ob_split_k(32, configured_split=2) == 2
+    with pytest.raises(ValueError, match="unsupported.*split-K factor"):
+        _select_dsv4_ob_split_k(32, configured_split=3)
+
+
+def test_dsv4_fmha_epilogue_output_uses_fused_oproj():
+    attn_fp8 = torch.empty((16, 4, 4096), device="meta", dtype=torch.float8_e4m3fn)
+    attn_scale = torch.empty((16, 32, 4), device="meta")
+    expected = torch.empty((8, 7168), device="meta", dtype=torch.bfloat16)
+    calls = []
+    module = SimpleNamespace(
+        _should_use_fused_oproj=lambda: True,
+        _fused_oa_ob_proj=lambda fp8, scale, num_tokens: calls.append((fp8, scale, num_tokens))
+        or expected,
+    )
+
+    output = MLA._deepseek_v4_o_proj(module, (attn_fp8, attn_scale))
+
+    assert output is expected
+    assert calls == [(attn_fp8, attn_scale, 4)]
+
+
+@pytest.mark.parametrize(
+    ("num_tokens", "num_splits", "expected"),
+    [
+        (1, 2, ((128, 32), (2, 1), True, None, True)),
+        (32, 2, ((128, 32), (2, 1), True, None, True)),
+        (64, 2, ((256, 64), (2, 1), True, 8, True)),
+        (96, 2, ((256, 128), (2, 1), True, None, True)),
+        (128, 2, ((256, 128), (2, 1), True, None, True)),
+        (256, 1, ((256, 128), (2, 1), True, None, True)),
+        (448, 1, ((256, 208), (2, 1), False, None, True)),
+        (512, 1, ((256, 208), (2, 1), False, None, True)),
+        (704, 1, ((256, 144), (2, 1), True, None, True)),
+        (2240, 1, ((256, 224), (2, 1), True, None, True)),
+        (2304, 1, ((256, 224), (2, 1), False, None, False)),
+        (3776, 1, ((256, 224), (2, 1), True, None, True)),
+        (4096, 1, ((256, 224), (2, 1), False, None, True)),
+        (5888, 1, ((256, 224), (2, 1), False, None, True)),
+        (9472, 1, ((256, 224), (2, 1), False, None, True)),
+        (16384, 1, ((256, 240), (2, 1), True, None, True)),
+    ],
+)
+def test_dsv4_ob_cute_tactic(num_tokens, num_splits, expected):
+    runner = cute_dsl_custom_ops.CuteDSLFp8SplitKGemmRunner
+    assert runner._get_tactic(num_tokens, 7168, num_splits, 74) == expected
+
+
+def test_dsv4_ob_split_k_one_uses_cute_dsl_and_caches_weight_scale(monkeypatch):
+    from tensorrt_llm import deep_gemm
+    from tensorrt_llm.quantization.utils import fp8_utils
+
+    m, n, k = 256, 7168, 16384
+    packed_k = k // 512
+    weight_scale_storage = torch.empty((packed_k, n), device="meta", dtype=torch.int32)
+    transformed_scale = torch.as_strided(weight_scale_storage, (n, packed_k), (1, n))
+    module = SimpleNamespace(
+        hidden_size=n,
+        dtype=torch.bfloat16,
+        ob_split_k=1,
+        o_b_proj=SimpleNamespace(
+            weight=torch.empty((n, k), device="meta", dtype=torch.float8_e4m3fn),
+            weight_scale=torch.empty((n // 128, k // 128), device="meta"),
+        ),
+    )
+    activation = torch.empty((m, k), device="meta", dtype=torch.float8_e4m3fn)
+    activation_scale = torch.empty_strided((m, packed_k), (1, m), device="meta", dtype=torch.int32)
+    calls = []
+    transforms = []
+
+    def fail_deep_gemm(*args, **kwargs):
+        raise AssertionError("supported SK1 must not dispatch to DeepGEMM")
+
+    def splitk_gemm(a, sfa, b, sfb, partials, num_splits):
+        calls.append((a, sfa, b, sfb, partials, num_splits))
+
+    def transform_weight_scale(scale, **kwargs):
+        transforms.append((scale, kwargs))
+        return transformed_scale
+
+    monkeypatch.setattr(attention_module, "IS_CUTLASS_DSL_AVAILABLE", True)
+    monkeypatch.setattr(deep_gemm, "fp8_gemm_nt", fail_deep_gemm)
+    monkeypatch.setattr(fp8_utils, "transform_sf_into_required_layout", transform_weight_scale)
+    monkeypatch.setattr(torch.ops.trtllm, "dsv4_fp8_splitk_gemm", splitk_gemm)
+    output = MLA._fused_ob_gemm(module, activation, activation_scale, m)
+    cached_output = MLA._fused_ob_gemm(module, activation, activation_scale, m)
+
+    assert output.shape == (m, n)
+    assert cached_output.shape == (m, n)
+    assert output.device.type == "meta"
+    assert len(transforms) == 1
+    assert transforms[0][0] is module.o_b_proj.weight_scale
+    assert module.o_b_proj._ob_wsf_int is transformed_scale
+    assert len(calls) == 2
+    assert calls[0][0] is activation
+    assert calls[0][1] is activation_scale
+    assert calls[0][2] is module.o_b_proj.weight
+    assert calls[0][3] is transformed_scale
+    assert calls[0][4].shape == (1, m, n)
+    assert calls[0][5] == 1
+
+
+@pytest.mark.parametrize("m", [1, 32, 64, 128])
+def test_dsv4_ob_auto_split_through_m128_uses_cute_dsl(m, monkeypatch):
+    from tensorrt_llm import deep_gemm
+
+    n, k = 7168, 16384
+    packed_k = k // 512
+    weight_scale_storage = torch.empty((packed_k, n), device="meta", dtype=torch.int32)
+    module = SimpleNamespace(
+        hidden_size=n,
+        dtype=torch.bfloat16,
+        ob_split_k=None,
+        o_b_proj=SimpleNamespace(
+            weight=torch.empty((n, k), device="meta", dtype=torch.float8_e4m3fn),
+            weight_scale=torch.as_strided(weight_scale_storage, (n, packed_k), (1, n)),
+        ),
+    )
+    activation = torch.empty((m, k), device="meta", dtype=torch.float8_e4m3fn)
+    aligned_m = (m + 3) // 4 * 4
+    activation_scale = torch.empty_strided(
+        (m, packed_k), (1, aligned_m), device="meta", dtype=torch.int32
+    )
+    calls = []
+
+    def fail_deep_gemm(*args, **kwargs):
+        raise AssertionError("automatic M<=128 SK2 must not dispatch to DeepGEMM")
+
+    def splitk_gemm(a, sfa, b, sfb, partials, num_splits):
+        calls.append((a, sfa, b, sfb, partials, num_splits))
+
+    monkeypatch.setattr(attention_module, "IS_CUTLASS_DSL_AVAILABLE", True)
+    monkeypatch.setattr(deep_gemm, "fp8_gemm_nt", fail_deep_gemm)
+    monkeypatch.setattr(torch.ops.trtllm, "dsv4_fp8_splitk_gemm", splitk_gemm)
+    output = MLA._fused_ob_gemm(module, activation, activation_scale, m)
+
+    assert output.shape == (2 * m, n)
+    assert len(calls) == 1
+    assert calls[0][4].shape == (2, m, n)
+    assert calls[0][5] == 2
+
+
+@pytest.mark.parametrize("m", [160])
+def test_dsv4_ob_unsplit_mid_m_uses_deep_gemm(m, monkeypatch):
+    from tensorrt_llm import deep_gemm
+    from tensorrt_llm.quantization.utils import fp8_utils
+
+    n, k = 7168, 16384
+    packed_k = k // 512
+    module = SimpleNamespace(
+        hidden_size=n,
+        dtype=torch.bfloat16,
+        ob_split_k=None,
+        o_b_proj=SimpleNamespace(
+            weight=torch.empty((n, k), device="meta", dtype=torch.float8_e4m3fn),
+            weight_scale=torch.empty((n // 128, k // 128), device="meta"),
+        ),
+    )
+    activation = torch.empty((m, k), device="meta", dtype=torch.float8_e4m3fn)
+    activation_scale = torch.empty_strided((m, packed_k), (1, m), device="meta", dtype=torch.int32)
+    calls = []
+
+    def deep_gemm_nt(inputs, weights, output, **kwargs):
+        calls.append((inputs, weights, output, kwargs))
+
+    def fail_cute(*args, **kwargs):
+        raise AssertionError("default unsplit M below 256 must dispatch to DeepGEMM")
+
+    def fail_transform(*args, **kwargs):
+        raise AssertionError("DeepGEMM fallback must not transform the weight scale")
+
+    monkeypatch.setattr(attention_module, "IS_CUTLASS_DSL_AVAILABLE", True)
+    monkeypatch.setattr(deep_gemm, "fp8_gemm_nt", deep_gemm_nt)
+    monkeypatch.setattr(fp8_utils, "transform_sf_into_required_layout", fail_transform)
+    monkeypatch.setattr(torch.ops.trtllm, "dsv4_fp8_splitk_gemm", fail_cute)
+    output = MLA._fused_ob_gemm(module, activation, activation_scale, m)
+
+    assert output.shape == (m, n)
+    assert output.device.type == "meta"
+    assert len(calls) == 1
+    assert calls[0][0] == (activation, activation_scale)
+    assert calls[0][1] == (module.o_b_proj.weight, module.o_b_proj.weight_scale)
+    assert calls[0][2] is output
+
+
+@skip_pre_blackwell
+@pytest.mark.parametrize(
+    ("num_tokens", "num_splits"),
+    [(1, 2), (2, 2), (32, 2), (64, 2), (128, 2), (256, 1), (512, 1), (64, 4)],
+)
+def test_dsv4_pro_fp8_splitk_gemm_partials(num_tokens: int, num_splits: int):
+    if get_sm_version() // 10 != 10:
+        pytest.skip("dsv4_fp8_splitk_gemm requires the SM100 family")
+
+    n, k = 7168, 16384
+    aligned_m = (num_tokens + 3) // 4 * 4
+    packed_k = k // 512
+    a = torch.full((num_tokens, k), 0.03125, device="cuda", dtype=torch.float8_e4m3fn)
+    b = torch.full((n, k), 0.03125, device="cuda", dtype=torch.float8_e4m3fn)
+
+    # Pack four unit UE8M0 scales per int32.
+    sfa_storage = torch.full((aligned_m * packed_k,), 0x7F7F7F7F, device="cuda", dtype=torch.int32)
+    sfb_storage = torch.full((n * packed_k,), 0x7F7F7F7F, device="cuda", dtype=torch.int32)
+    sfa = torch.as_strided(sfa_storage, (num_tokens, packed_k), (1, aligned_m))
+    sfb = torch.as_strided(sfb_storage, (n, packed_k), (1, n))
+    partials = torch.empty((num_splits, num_tokens, n), device="cuda", dtype=torch.bfloat16)
+
+    torch.ops.trtllm.dsv4_fp8_splitk_gemm(a, sfa, b, sfb, partials, num_splits)
+    expected_partial = (k // num_splits) * 0.03125 * 0.03125
+    torch.testing.assert_close(
+        partials.float(), torch.full_like(partials, expected_partial, dtype=torch.float32)
+    )
+
+
+@skip_pre_blackwell
+def test_dsv4_pro_fp8_splitk_gemm_packed_scales():
+    if get_sm_version() // 10 != 10:
+        pytest.skip("dsv4_fp8_splitk_gemm requires the SM100 family")
+
+    m, n, k, num_splits = 64, 128, 2048, 4
+    k_blocks = k // 128
+    torch.manual_seed(1234)
+    a = (torch.randn((m, k), device="cuda") * 0.125).to(torch.float8_e4m3fn)
+    b = (torch.randn((n, k), device="cuda") * 0.125).to(torch.float8_e4m3fn)
+    exp_a = torch.randint(124, 131, (m, k_blocks), device="cuda")
+    exp_b = torch.randint(124, 131, (n // 128, k_blocks), device="cuda").repeat_interleave(
+        128, dim=0
+    )
+
+    def pack_scales(exponents: torch.Tensor, aligned_rows: int) -> torch.Tensor:
+        rows, num_blocks = exponents.shape
+        grouped = exponents.reshape(rows, num_blocks // 4, 4).to(torch.int32)
+        packed = (
+            grouped[..., 0]
+            | (grouped[..., 1] << 8)
+            | (grouped[..., 2] << 16)
+            | (grouped[..., 3] << 24)
+        )
+        storage = torch.zeros((num_blocks // 4, aligned_rows), device="cuda", dtype=torch.int32)
+        storage[:, :rows] = packed.transpose(0, 1)
+        return torch.as_strided(storage, packed.shape, (1, aligned_rows))
+
+    sfa = pack_scales(exp_a, m)
+    sfb = pack_scales(exp_b, n)
+    partials = torch.empty((num_splits, m, n), device="cuda", dtype=torch.bfloat16)
+    torch.ops.trtllm.dsv4_fp8_splitk_gemm(a, sfa, b, sfb, partials, num_splits)
+
+    scale_a = torch.exp2(exp_a.float() - 127.0)
+    scale_b = torch.exp2(exp_b.float() - 127.0)
+    a_dequant = (a.float().reshape(m, k_blocks, 128) * scale_a[..., None]).reshape(m, k)
+    b_dequant = (b.float().reshape(n, k_blocks, 128) * scale_b[..., None]).reshape(n, k)
+    split_k = k // num_splits
+    expected = torch.stack(
+        [
+            a_dequant[:, split * split_k : (split + 1) * split_k]
+            @ b_dequant[:, split * split_k : (split + 1) * split_k].T
+            for split in range(num_splits)
+        ]
+    )
+    torch.testing.assert_close(partials.float(), expected, rtol=0.02, atol=0.05)
+
+
 def _per_token_fp8_quant_dequant(x: torch.Tensor) -> torch.Tensor:
     """Simulate the fused inverse-RoPE FP8 quantization consumed by o_a_proj."""
     original_shape = x.shape
@@ -90,20 +366,16 @@ def calculate_reference_deepseek_v4_o_proj(
     return output
 
 
-@skip_pre_blackwell
-@pytest.mark.skip_less_device_memory(80000)
-@pytest.mark.parametrize("num_tokens", [1, 16, 128])
-@pytest.mark.parametrize("dtype_str", ["bf16", "fp8"])
-def test_deepseek_v4_o_proj(num_tokens: int, dtype_str: str):
-    """Test DeepSeek-V4 output projection (_deepseek_v4_o_proj)."""
-    print(
-        f"\n{'=' * 80}\nTesting: deepseek_v4_o_proj num_tokens={num_tokens} dtype={dtype_str}\n{'=' * 80}"
-    )
+def _build_dsv4_o_proj_case(num_tokens: int, dtype_str: str, device: torch.device):
+    """Build an MLA module, inputs, and reference-path tensors for the DeepSeek-V4
+    o_proj tests. Shared by the correctness test and the default fused
+    fp8-equivalence test so both exercise an identical setup.
 
-    if dtype_str == "fp8" and get_sm_version() < 100:
-        pytest.skip("FP8 is not supported on pre-Blackwell architectures")
-
-    device = torch.device("cuda")
+    Returns:
+        (mla, attn_out_latent, position_ids, refs) where ``refs`` is a namespace
+        carrying the dequantized weights / freqs_cis / dims the analytic
+        reference path consumes.
+    """
     dtype = torch.bfloat16
 
     # Model configuration matching the reference model
@@ -197,6 +469,7 @@ def test_deepseek_v4_o_proj(num_tokens: int, dtype_str: str):
 
     # Initialize weights
     nn_init_std = 0.02
+    fp8_a_weight = fp8_a_scale = o_a_proj_bf16 = fp8_b_weight_dequant = dim = None
     with torch.no_grad():
         # Initialize o_a_proj weights
         if dtype_str == "bf16":
@@ -221,12 +494,7 @@ def test_deepseek_v4_o_proj(num_tokens: int, dtype_str: str):
             fp8_a_weight = fp8_a_weight.reshape(n_local_groups, o_lora_rank, dim)
             mla.o_a_proj.data = fp8_a_weight
             mla.o_a_proj_scale.data = fp8_a_scale
-            # mla.o_a_proj_dequant is None for DSv4 on SM100: PR #14254
-            # decouples the FP8-native o_a_proj path from
-            # use_cute_dsl_blockscaling_bmm, so DSv4 unconditionally uses the
-            # fused inv-RoPE + FP8 quant + cute-dsl BMM chain and never needs
-            # the bf16-dequant fallback buffer. The reference path below uses
-            # o_a_proj_bf16 directly.
+            # SM100 keeps only the quantized O_a weights.
 
         # Initialize o_b_proj weights
         if dtype_str == "bf16":
@@ -254,17 +522,11 @@ def test_deepseek_v4_o_proj(num_tokens: int, dtype_str: str):
     attn_out_latent = torch.randn(num_tokens, num_heads, qk_head_dim, dtype=dtype, device=device)
     position_ids = torch.arange(num_tokens, dtype=torch.int32, device=device)
 
-    # Call the deepseek_v4 output projection (mla_rope_inplace modifies attn_out_latent
-    # in-place, so clone before passing to preserve original for reference)
-    output = mla._deepseek_v4_o_proj(attn_out_latent.clone(), position_ids)
-
-    # Calculate reference output
+    # Dequantize the weights actually consumed by the FP8 path.
     if dtype_str == "bf16":
         o_a_proj_ref = mla.o_a_proj.data
         o_b_proj_weight_ref = mla.o_b_proj.weight.data
     else:
-        # Match the FP8-native o_a_proj path: the runtime BMM consumes
-        # quantized o_a_proj plus block scales, not the original BF16 weight.
         o_a_proj_ref = (
             weight_dequant(
                 fp8_a_weight.reshape(-1, dim).contiguous(),
@@ -285,14 +547,47 @@ def test_deepseek_v4_o_proj(num_tokens: int, dtype_str: str):
         rope_config.rope_scaling["beta_slow"],
     ).to(device)
 
-    reference_output = calculate_reference_deepseek_v4_o_proj(
-        attn_out_latent=attn_out_latent,
+    refs = SimpleNamespace(
         o_a_proj=o_a_proj_ref,
         o_b_proj_weight=o_b_proj_weight_ref,
-        freqs_cis=freqs_cis[0:num_tokens],
+        freqs_cis=freqs_cis,
         n_local_groups=n_local_groups,
         qk_nope_head_dim=qk_nope_head_dim,
         qk_rope_head_dim=qk_rope_head_dim,
+    )
+    return mla, attn_out_latent, position_ids, refs
+
+
+@skip_pre_blackwell
+@pytest.mark.skip_less_device_memory(80000)
+@pytest.mark.parametrize("num_tokens", [1, 16, 128])
+@pytest.mark.parametrize("dtype_str", ["bf16", "fp8"])
+def test_deepseek_v4_o_proj(num_tokens: int, dtype_str: str):
+    """Test DeepSeek-V4 output projection (_deepseek_v4_o_proj)."""
+    print(
+        f"\n{'=' * 80}\nTesting: deepseek_v4_o_proj num_tokens={num_tokens} dtype={dtype_str}\n{'=' * 80}"
+    )
+
+    if dtype_str == "fp8" and get_sm_version() < 100:
+        pytest.skip("FP8 is not supported on pre-Blackwell architectures")
+
+    device = torch.device("cuda")
+
+    mla, attn_out_latent, position_ids, refs = _build_dsv4_o_proj_case(
+        num_tokens, dtype_str, device
+    )
+
+    # Preserve the input because inverse RoPE is in-place.
+    output = mla._deepseek_v4_o_proj(attn_out_latent.clone(), position_ids)
+
+    reference_output = calculate_reference_deepseek_v4_o_proj(
+        attn_out_latent=attn_out_latent,
+        o_a_proj=refs.o_a_proj,
+        o_b_proj_weight=refs.o_b_proj_weight,
+        freqs_cis=refs.freqs_cis[0:num_tokens],
+        n_local_groups=refs.n_local_groups,
+        qk_nope_head_dim=refs.qk_nope_head_dim,
+        qk_rope_head_dim=refs.qk_rope_head_dim,
         device=device,
         is_fp8=dtype_str == "fp8",
     )
@@ -321,3 +616,91 @@ def test_deepseek_v4_o_proj(num_tokens: int, dtype_str: str):
     else:
         torch.testing.assert_close(output, reference_output, rtol=0.1, atol=0.1)
         print(f"  ✓ Test passed for num_tokens={num_tokens}, dtype={dtype_str}\n")
+
+
+@skip_pre_blackwell
+@pytest.mark.skip_less_device_memory(80000)
+@pytest.mark.parametrize("num_tokens", [1, 16, 32, 128, 256])
+def test_deepseek_v4_o_proj_fused_fp8_equivalence(num_tokens: int, monkeypatch):
+    """The default fused FP8 epilogue must match the disabled baseline.
+
+    Fused (default): o_a's CuTe-DSL GEMM emits o_lora directly as fp8
+    e4m3 + packed-UE8M0 1x128 scale factors, fed straight to o_b.
+    Disabled: o_a emits bf16 o_lora, then ``o_b_proj`` runs the separate 1x128
+    quant + DeepGEMM. Since the fusion only folds the *same* quant into o_a's
+    epilogue, the two paths must match far tighter than the fp8-vs-bf16 reference.
+    """
+    if get_sm_version() < 100:
+        pytest.skip("fused DeepSeek-V4 FP8 O-projection requires Blackwell (SM100+)")
+
+    device = torch.device("cuda")
+    mla, attn_out_latent, position_ids, refs = _build_dsv4_o_proj_case(num_tokens, "fp8", device)
+
+    # Ensure the fused path is eligible.
+    assert mla.o_a_proj.dtype == torch.float8_e4m3fn
+    assert mla.n_local_groups == mla.num_groups
+    assert getattr(mla.o_b_proj, "tp_size", 1) == 1
+    assert mla.o_b_proj.has_fp8_block_scales
+    assert not getattr(mla.o_b_proj, "use_cute_dsl_blockscaling_mm", False)
+
+    # Explicit kill switch retains the unfused fallback.
+    monkeypatch.setenv("TRTLLM_DSV4_DISABLE_FUSED_OPROJ", "1")
+    out_unfused = mla._deepseek_v4_o_proj(attn_out_latent.clone(), position_ids)
+
+    # Unset is the production default and must take the fused path.
+    monkeypatch.delenv("TRTLLM_DSV4_DISABLE_FUSED_OPROJ", raising=False)
+    mla.ob_split_k = 1
+    out_fused = mla._deepseek_v4_o_proj(attn_out_latent.clone(), position_ids)
+
+    assert out_fused.shape == out_unfused.shape, (
+        f"Shape mismatch: fused {out_fused.shape} vs unfused {out_unfused.shape}"
+    )
+    assert out_fused.dtype == out_unfused.dtype, (
+        f"Dtype mismatch: fused {out_fused.dtype} vs unfused {out_unfused.dtype}"
+    )
+    assert torch.isfinite(out_fused).all(), "Fused output contains non-finite values"
+
+    # Analytic reference (same correctness bar as the unfused correctness test).
+    reference_output = calculate_reference_deepseek_v4_o_proj(
+        attn_out_latent=attn_out_latent.clone(),
+        o_a_proj=refs.o_a_proj,
+        o_b_proj_weight=refs.o_b_proj_weight,
+        freqs_cis=refs.freqs_cis[0:num_tokens],
+        n_local_groups=refs.n_local_groups,
+        qk_nope_head_dim=refs.qk_nope_head_dim,
+        qk_rope_head_dim=refs.qk_rope_head_dim,
+        device=device,
+        is_fp8=True,
+    )
+
+    diff_vs_ref = _calc_diff(out_fused, reference_output)
+    diff_vs_unfused = _calc_diff(out_fused, out_unfused)
+    print(
+        f"\n  num_tokens={num_tokens}  diff(fused,ref)={diff_vs_ref:.3e}  "
+        f"diff(fused,unfused)={diff_vs_unfused:.3e}"
+    )
+
+    assert diff_vs_ref < FP8_O_PROJ_DIFF_TOL, f"fused vs reference {diff_vs_ref=}"
+    assert diff_vs_unfused < 1e-3, f"fused vs unfused {diff_vs_unfused=}"
+
+    expected_smem_epilogue = num_tokens <= 32
+    expected_smem_row_iters = (num_tokens + 15) // 16 if expected_smem_epilogue else 1
+    fp8out_keys = [
+        key
+        for key in cute_dsl_custom_ops.CuteDSLFp8BlackwellBmmRunner.kernel_cache
+        if key[0] == "fp8out"
+        and key[-2] == expected_smem_row_iters
+        and key[-1] == expected_smem_epilogue
+    ]
+    assert len(fp8out_keys) == 1
+
+    if num_tokens == 16:
+        compiled_gemm = cute_dsl_custom_ops.CuteDSLFp8BlackwellBmmRunner.kernel_cache[
+            fp8out_keys[0]
+        ]
+        out_cached = mla._deepseek_v4_o_proj(attn_out_latent.clone(), position_ids)
+        assert (
+            cute_dsl_custom_ops.CuteDSLFp8BlackwellBmmRunner.kernel_cache[fp8out_keys[0]]
+            is compiled_gemm
+        )
+        torch.testing.assert_close(out_cached, out_unfused, rtol=0, atol=0)
diff --git a/tests/unittest/_torch/modules/test_mhc.py b/tests/unittest/_torch/modules/test_mhc.py
index cf03e050ad61..534fd0d7473e 100644
--- a/tests/unittest/_torch/modules/test_mhc.py
+++ b/tests/unittest/_torch/modules/test_mhc.py
@@ -782,6 +782,61 @@ def make_runner_inputs():
         )
 
 
+@pytest.mark.parametrize(
+    "tactic",
+    [
+        ("fused_half_mma", 0, 0, 0, 1),
+        ("fused_half_fma", 3, 2, 256, 1),
+    ],
+)
+def test_mhc_fused_hc_reduces_split_x_for_production_backends(tactic):
+    """Production half backends consume split-major O_b partials correctly."""
+    from tensorrt_llm._torch.modules.mhc.mhc_cuda import MhcFusedHcRunner
+
+    n, hidden_size, hc_mult, x_splits = 32, 7168, 4, 4
+    pre_data = generate_realistic_pre_data(n=n, hc_mult=hc_mult, hidden_size=hidden_size)
+    torch.random.manual_seed(23)
+
+    x_partials = (
+        torch.randn((x_splits, n, hidden_size), dtype=torch.float32, device="cuda") * 0.25
+    ).bfloat16()
+    x_reduced = x_partials.float().sum(dim=0).bfloat16()
+    residual_prev = torch.randn(
+        (n, hc_mult, hidden_size), dtype=torch.float32, device="cuda"
+    ).bfloat16()
+    post_mix_prev = torch.randn((n, hc_mult), dtype=torch.float32, device="cuda") * 0.1
+    comb_mix_prev = torch.randn((n, hc_mult, hc_mult), dtype=torch.float32, device="cuda") * 0.1
+
+    runner = MhcFusedHcRunner(
+        n=hc_mult,
+        hidden_size=hidden_size,
+        rms_eps=pre_data["rms_eps"],
+        hc_pre_eps=pre_data["hc_pre_eps"],
+        hc_sinkhorn_eps=pre_data["hc_sinkhorn_eps"],
+        hc_post_mult_value=pre_data["hc_post_mult_value"],
+        sinkhorn_repeat=pre_data["sinkhorn_repeat"],
+    )
+
+    def inputs(x_prev):
+        return [
+            x_prev.contiguous(),
+            residual_prev,
+            post_mix_prev,
+            comb_mix_prev,
+            pre_data["fn"].contiguous(),
+            pre_data["hc_scale"].contiguous(),
+            pre_data["hc_base"].contiguous(),
+        ]
+
+    reference = tuple(output.clone() for output in runner(inputs(x_reduced), tactic=tactic))
+    split_output = runner(inputs(x_partials.reshape(x_splits * n, hidden_size)), tactic=tactic)
+
+    torch.testing.assert_close(split_output[0], reference[0], rtol=1e-2, atol=2e-2)
+    torch.testing.assert_close(split_output[1], reference[1], rtol=3e-3, atol=5e-3)
+    torch.testing.assert_close(split_output[2], reference[2], rtol=3e-3, atol=5e-3)
+    torch.testing.assert_close(split_output[3], reference[3], rtol=1e-2, atol=2e-2)
+
+
 @pytest.mark.parametrize(
     "tactic",
     [