diff --git a/benchmark/gemm/frost/benchmark_block_scale_matmul.py b/benchmark/gemm/frost/benchmark_block_scale_matmul.py
index d277f919a..3a444ed95 100644
--- a/benchmark/gemm/frost/benchmark_block_scale_matmul.py
+++ b/benchmark/gemm/frost/benchmark_block_scale_matmul.py
@@ -25,6 +25,7 @@
 
 from cudnn.gemm.frost.compiler import jit_from_cudnn_graph
 from cudnn.gemm.frost.tile_config import CATALOG as _CATALOG
+from cudnn.gemm.frost.tile_config import by_name as _by_name
 
 
 def _build_spec_map():
@@ -36,8 +37,8 @@ def _build_spec_map():
         kb_want = 384 if cfg.pipeline == "sm103" else 128
         if cfg.cta_tile_m % 128 or cfg.cta_tile_n % 128 or cfg.cta_tile_k_bytes != kb_want:
             continue
-        # sm103 has 1ctamma + 2ctamma CLC templates (no static variants).
-        scheds = (("clc", ""),) if cfg.pipeline == "sm103" else (("clc", ""), ("static", "_static"))
+        # Only sm100 has static-scheduler variants; sm103 / sm107 are CLC-only.
+        scheds = (("clc", ""), ("static", "_static")) if cfg.pipeline == "sm100" else (("clc", ""),)
         for cg in (1, 2):
             if cg == 2 and (cfg.cgrp_size_m % 2 or cfg.cta_tile_m == 64):
                 continue
@@ -48,6 +49,27 @@ def _build_spec_map():
 
 _SPEC_MAP = _build_spec_map()
 
+_LABEL_RE = re.compile(r"^(CONFIG_sm\d+_\d+x\d+x\d+_\d+x\d+x\d+_cluster\d+x\d+)_([12])ctamma(_static)?$")
+
+
+def _spec_for(name):
+    """(geometry cfg, cta_group, scheduler) for a --configs label, or None.
+
+    The sweep set comes from the registry funnel over CATALOG; a label naming a
+    geometry outside it (e.g. a num_mma_m > 1 tile, which `by_name` synthesizes) is
+    still runnable, so parse it rather than reporting UNKNOWN_CONFIG."""
+    spec = _SPEC_MAP.get(name)
+    if spec is not None:
+        return spec
+    m = _LABEL_RE.match(name)
+    if m is None:
+        return None
+    try:
+        cfg = _by_name(m.group(1))
+    except (KeyError, NotImplementedError):
+        return None
+    return cfg, int(m.group(2)), "static" if m.group(3) else "clc"
+
 
 def _vp_bs(handles, a, b, c, sfa, sfb):
     """Block-scale single-GEMM variant-pack dict keyed by the graph's tensors."""
@@ -57,7 +79,7 @@ def _vp_bs(handles, a, b, c, sfa, sfb):
 
 def _build_plan(g, cfg, name):
     """JIT-compile the recorded graph with a forced tile config."""
-    return jit_from_cudnn_graph(g, config=cfg, cta_group=_SPEC_MAP[name][1], scheduler=_SPEC_MAP[name][2])
+    return jit_from_cudnn_graph(g, config=cfg, cta_group=_spec_for(name)[1], scheduler=_spec_for(name)[2])
 
 
 # Combo table (input dtype family + scale dtype + block size)
@@ -508,10 +530,10 @@ def _nsys_worker(shape, combo, configs, warmup, iters, ref_mode, nbuf) -> None:
     torch.cuda.synchronize()
 
     # 2. each block-scale config.
-    name_to_cfg = {lbl: sp[0] for lbl, sp in _SPEC_MAP.items()}
     config_names = configs or list(_SPEC_MAP)
     for name in config_names:
-        cfg = name_to_cfg.get(name)
+        spec = _spec_for(name)
+        cfg = spec[0] if spec else None
         if cfg is None:
             continue
         try:
@@ -603,7 +625,6 @@ def main() -> int:
 
     flops = 2 * B * M * N * K
     config_names = [c.strip() for c in args.configs.split(",")] if args.configs else list(_SPEC_MAP)
-    name_to_cfg = {lbl: sp[0] for lbl, sp in _SPEC_MAP.items()}
 
     print(f"\n=== block-scale matmul B={B} {M}x{N}x{K}  (~{flops / 1e9:.1f} GFLOP) — " f"{combo} in / BF16 out ===")
 
@@ -644,7 +665,8 @@ def _fmt_row(name: str, tflops: float, ms: float, note: str, ref_tflops: float)
             print("  reference kernel: not detected in nsys output")
 
         for name in config_names:
-            cfg = name_to_cfg.get(name)
+            spec = _spec_for(name)
+            cfg = spec[0] if spec else None
             if cfg is None:
                 rows.append((name, 0.0, float("inf"), "UNKNOWN_CONFIG"))
                 continue
@@ -687,7 +709,8 @@ def _fmt_row(name: str, tflops: float, ms: float, note: str, ref_tflops: float)
 
         ctx_dead = False
         for name in config_names:
-            cfg = name_to_cfg.get(name)
+            spec = _spec_for(name)
+            cfg = spec[0] if spec else None
             if cfg is None:
                 row = (name, 0.0, float("inf"), "UNKNOWN_CONFIG")
             elif ctx_dead:
diff --git a/benchmark/gemm/frost/benchmark_block_scale_matmul_swiglu.py b/benchmark/gemm/frost/benchmark_block_scale_matmul_swiglu.py
index f90f39263..59d434955 100644
--- a/benchmark/gemm/frost/benchmark_block_scale_matmul_swiglu.py
+++ b/benchmark/gemm/frost/benchmark_block_scale_matmul_swiglu.py
@@ -11,6 +11,7 @@
 from __future__ import annotations
 
 import argparse
+import re
 import sys
 from typing import Callable
 
@@ -21,6 +22,7 @@
 from types import SimpleNamespace
 
 from cudnn.gemm.frost.compiler import jit_from_cudnn_graph
+from cudnn.gemm.frost.tile_config import by_name as _by_name
 from cudnn.gemm.frost.graph_analyzer import analyze
 from cudnn.gemm.frost.kernel_registry import candidates as _registry_candidates
 
@@ -204,6 +206,27 @@ def _build_spec_map():
 
 _SPEC_MAP = _build_spec_map()
 
+_LABEL_RE = re.compile(r"^(CONFIG_sm\d+_\d+x\d+x\d+_\d+x\d+x\d+_cluster\d+x\d+)_([12])ctamma(_static)?$")
+
+
+def _spec_for(name):
+    """(geometry cfg, cta_group, scheduler) for a --configs label, or None.
+
+    The sweep set comes from the registry funnel over CATALOG; a label naming a
+    geometry outside it (e.g. a num_mma_m > 1 tile, which `by_name` synthesizes) is
+    still runnable, so parse it rather than reporting it unsweepable."""
+    spec = _SPEC_MAP.get(name)
+    if spec is not None:
+        return spec
+    m = _LABEL_RE.match(name)
+    if m is None:
+        return None
+    try:
+        cfg = _by_name(m.group(1))
+    except (KeyError, NotImplementedError):
+        return None
+    return cfg, int(m.group(2)), "static" if m.group(3) else "clc"
+
 
 def main() -> int:
     p = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter)
@@ -244,10 +267,11 @@ def main() -> int:
 
     best = None
     for name in config_names:
-        if name not in _SPEC_MAP:
+        spec = _spec_for(name)
+        if spec is None:
             print(f"  {name:62s} UNKNOWN (not a sweepable block-scale strategy)")
             continue
-        cfg, cta_group, sched = _SPEC_MAP[name]
+        cfg, cta_group, sched = spec
         try:
             g, h = _graph(B, M, N, K)
             plan = _build_plan(g, cfg, cta_group, sched)
diff --git a/benchmark/gemm/frost/benchmark_matmul_mixed_input.py b/benchmark/gemm/frost/benchmark_matmul_mixed_input.py
index 45faa5020..3c9d45a3e 100644
--- a/benchmark/gemm/frost/benchmark_matmul_mixed_input.py
+++ b/benchmark/gemm/frost/benchmark_matmul_mixed_input.py
@@ -27,6 +27,7 @@
 import torch
 
 from cudnn.gemm.frost.compiler import jit_from_cudnn_graph
+from cudnn.gemm.frost.tile_config import by_name as _by_name
 from cudnn.gemm.frost.graph_analyzer import analyze
 from cudnn.gemm.frost.kernel_registry import candidates as _candidates
 
@@ -81,6 +82,27 @@ def _build_spec_map():
 
 _SPEC_MAP = _build_spec_map()
 
+_LABEL_RE = re.compile(r"^(CONFIG_sm\d+_\d+x\d+x\d+_\d+x\d+x\d+_cluster\d+x\d+)_([12])ctamma(_static)?$")
+
+
+def _spec_for(name):
+    """(geometry cfg, cta_group, scheduler) for a --configs label, or None.
+
+    The sweep set comes from the registry funnel over CATALOG; a label naming a
+    geometry outside it (e.g. a num_mma_m > 1 tile, which `by_name` synthesizes) is
+    still runnable, so parse it rather than reporting UNKNOWN_CONFIG."""
+    spec = _SPEC_MAP.get(name)
+    if spec is not None:
+        return spec
+    m = _LABEL_RE.match(name)
+    if m is None:
+        return None
+    try:
+        cfg = _by_name(m.group(1))
+    except (KeyError, NotImplementedError):
+        return None
+    return cfg, int(m.group(2)), "static" if m.group(3) else "clc"
+
 
 def _vp(handles, a, b, c):
     """Variant-pack dict {tensor: buffer}; `a` is the narrow (load-dtype) A root operand."""
@@ -90,7 +112,7 @@ def _vp(handles, a, b, c):
 
 def _build_plan(g, cfg, name):
     """JIT-compile the graph with a forced tile config -> callable kernel."""
-    return jit_from_cudnn_graph(g, config=cfg, cta_group=_SPEC_MAP[name][1], scheduler=_SPEC_MAP[name][2])
+    return jit_from_cudnn_graph(g, config=cfg, cta_group=_spec_for(name)[1], scheduler=_spec_for(name)[2])
 
 
 # ---------------------------------------------------------------------------
@@ -378,10 +400,10 @@ def _nsys_worker(shape, configs, warmup, iters, nbuf, load_dt, tin_dt, tout_dt)
     torch.cuda.synchronize()
 
     # 2. each GEMM config.
-    name_to_cfg = {lbl: sp[0] for lbl, sp in _SPEC_MAP.items()}
     config_names = configs or list(_SPEC_MAP)
     for name in config_names:
-        cfg = name_to_cfg.get(name)
+        spec = _spec_for(name)
+        cfg = spec[0] if spec else None
         if cfg is None:
             continue
         try:
@@ -449,7 +471,6 @@ def main() -> int:
 
     flops = 2 * B * M * N * K
     config_names = [c.strip() for c in args.configs.split(",")] if args.configs else list(_SPEC_MAP)
-    name_to_cfg = {lbl: sp[0] for lbl, sp in _SPEC_MAP.items()}
 
     print(f"\n=== mixed-input matmul B={B} {M}x{N}x{K}  (~{flops / 1e9:.1f} GFLOP) " f"— A={load_dt} -> {tin_dt} @ {tin_dt} -> {tout_dt} ===")
 
@@ -493,7 +514,8 @@ def _fmt_row(name, tflops, ms, note, ref_tflops) -> str:
             print("  cuBLAS kernel: not detected in nsys output")
 
         for name in config_names:
-            cfg = name_to_cfg.get(name)
+            spec = _spec_for(name)
+            cfg = spec[0] if spec else None
             if cfg is None:
                 rows.append((name, 0.0, float("inf"), "UNKNOWN_CONFIG"))
                 continue
@@ -528,7 +550,8 @@ def _fmt_row(name, tflops, ms, note, ref_tflops) -> str:
 
         ctx_dead = False
         for name in config_names:
-            cfg = name_to_cfg.get(name)
+            spec = _spec_for(name)
+            cfg = spec[0] if spec else None
             if cfg is None:
                 row = (name, 0.0, float("inf"), "UNKNOWN_CONFIG")
             elif ctx_dead:
diff --git a/benchmark/gemm/frost/benchmark_matmul_swiglu.py b/benchmark/gemm/frost/benchmark_matmul_swiglu.py
index ec6a72199..5f9ffb36e 100644
--- a/benchmark/gemm/frost/benchmark_matmul_swiglu.py
+++ b/benchmark/gemm/frost/benchmark_matmul_swiglu.py
@@ -10,6 +10,7 @@
 from __future__ import annotations
 
 import argparse
+import re
 import sys
 import time
 from typing import Callable
@@ -21,6 +22,7 @@
 from types import SimpleNamespace
 
 from cudnn.gemm.frost.compiler import jit_from_cudnn_graph
+from cudnn.gemm.frost.tile_config import by_name as _by_name
 from cudnn.gemm.frost.graph_analyzer import analyze
 from cudnn.gemm.frost.kernel_registry import candidates as _registry_candidates
 
@@ -147,6 +149,27 @@ def _build_spec_map():
 
 _SPEC_MAP = _build_spec_map()
 
+_LABEL_RE = re.compile(r"^(CONFIG_sm\d+_\d+x\d+x\d+_\d+x\d+x\d+_cluster\d+x\d+)_([12])ctamma(_static)?$")
+
+
+def _spec_for(name):
+    """(geometry cfg, cta_group, scheduler) for a --configs label, or None.
+
+    The sweep set comes from the registry funnel over CATALOG; a label naming a
+    geometry outside it (e.g. a num_mma_m > 1 tile, which `by_name` synthesizes) is
+    still runnable, so parse it rather than reporting it unsweepable."""
+    spec = _SPEC_MAP.get(name)
+    if spec is not None:
+        return spec
+    m = _LABEL_RE.match(name)
+    if m is None:
+        return None
+    try:
+        cfg = _by_name(m.group(1))
+    except (KeyError, NotImplementedError):
+        return None
+    return cfg, int(m.group(2)), "static" if m.group(3) else "clc"
+
 
 # ---------------------------------------------------------------------------
 # Main
@@ -204,10 +227,11 @@ def main() -> int:
 
     best = None
     for label in config_names:
-        if label not in _SPEC_MAP:
+        spec = _spec_for(label)
+        if spec is None:
             print(f"  {label:62s} UNKNOWN (not a sweepable swiglu strategy)")
             continue
-        cfg, cta_group, sched = _SPEC_MAP[label]
+        cfg, cta_group, sched = spec
         try:
             g, h = _graph_swiglu(B, M, N, K, in_dt, out_dt)
             plan = _build_plan(g, cfg, cta_group, sched)
diff --git a/benchmark/gemm/frost/benchmark_moe_block_scale_matmul.py b/benchmark/gemm/frost/benchmark_moe_block_scale_matmul.py
index 28fe855e6..a534171ec 100644
--- a/benchmark/gemm/frost/benchmark_moe_block_scale_matmul.py
+++ b/benchmark/gemm/frost/benchmark_moe_block_scale_matmul.py
@@ -25,6 +25,7 @@
 import torch
 
 from cudnn.gemm.frost.compiler import jit_from_cudnn_graph
+from cudnn.gemm.frost.tile_config import by_name as _by_name
 from cudnn.gemm.frost.graph_analyzer import analyze
 from cudnn.gemm.frost.kernel_registry import candidates as _candidates
 from cudnn.gemm.frost.tile_config import TileConfig
@@ -38,7 +39,7 @@ def _vp_moe_bs(handles, token, weight, sfa, sfb, fto, output):
 
 def _build_plan(g, cfg, name):
     """JIT-compile the recorded graph with a forced tile config -> compiled kernel."""
-    return jit_from_cudnn_graph(g, config=cfg, cta_group=_SPEC_MAP[name][1], scheduler=_SPEC_MAP[name][2])
+    return jit_from_cudnn_graph(g, config=cfg, cta_group=_spec_for(name)[1], scheduler=_spec_for(name)[2])
 
 
 # combo : (is_fp4, block_size, a_dtype, sf_dtype)
@@ -109,6 +110,27 @@ def _build_spec_map():
 
 _SPEC_MAP = _build_spec_map()
 
+_LABEL_RE = re.compile(r"^(CONFIG_sm\d+_\d+x\d+x\d+_\d+x\d+x\d+_cluster\d+x\d+)_([12])ctamma(_static)?$")
+
+
+def _spec_for(name):
+    """(geometry cfg, cta_group, scheduler) for a --configs label, or None.
+
+    The sweep set comes from the registry funnel over CATALOG; a label naming a
+    geometry outside it (e.g. a num_mma_m > 1 tile, which `by_name` synthesizes) is
+    still runnable, so parse it rather than reporting UNKNOWN_CONFIG."""
+    spec = _SPEC_MAP.get(name)
+    if spec is not None:
+        return spec
+    m = _LABEL_RE.match(name)
+    if m is None:
+        return None
+    try:
+        cfg = _by_name(m.group(1))
+    except (KeyError, NotImplementedError):
+        return None
+    return cfg, int(m.group(2)), "static" if m.group(3) else "clc"
+
 
 def _offsets(S: int, E: int) -> torch.Tensor:
     """Even split: group g owns rows [g*group_m, (g+1)*group_m)."""
@@ -425,9 +447,9 @@ def _nsys_worker(shape, combo, configs, warmup, iters, nbuf, no_baseline=False)
     pool = _mkdata_pool(S, N, K, E, combo, nbuf)
 
     # 2. each MoE-block-scale config.
-    name_to_cfg = {lbl: sp[0] for lbl, sp in _SPEC_MAP.items()}
     for name in configs or list(_SPEC_MAP):
-        cfg = name_to_cfg.get(name)
+        spec = _spec_for(name)
+        cfg = spec[0] if spec else None
         if cfg is None:
             continue
         try:
@@ -501,7 +523,6 @@ def main() -> int:
 
     flops = 2 * S * N * K
     config_names = [c.strip() for c in args.configs.split(",")] if args.configs else list(_SPEC_MAP)
-    name_to_cfg = {lbl: sp[0] for lbl, sp in _SPEC_MAP.items()}
 
     print(f"\n=== moe_block_scale_matmul G={G} M={M} N={N} K={K}  " f"(S={S} tokens, ~{flops / 1e9:.1f} GFLOP) — {combo} ===")
 
@@ -536,11 +557,12 @@ def _fmt_row(name, tflops, ms, note, ref_tflops) -> str:
             cublas_tflops, cublas_ms = float("nan"), float("nan")
             print("  cuBLAS kernel: not detected in nsys output")
         for name in config_names:
-            cfg = name_to_cfg.get(name)
+            spec = _spec_for(name)
+            cfg = spec[0] if spec else None
             if cfg is None:
                 rows.append((name, 0.0, float("inf"), "UNKNOWN_CONFIG"))
                 continue
-            tok = _kernel_match_token(cfg, _SPEC_MAP[name][1])
+            tok = _kernel_match_token(cfg, _spec_for(name)[1])
             matches = [(k, v) for k, v in kern_times.items() if tok in k]
             if not matches:
                 rows.append((name, 0.0, float("inf"), "NO_KERNEL_IN_NSYS"))
@@ -587,7 +609,8 @@ def _fmt_row(name, tflops, ms, note, ref_tflops) -> str:
 
         ctx_dead = False
         for name in config_names:
-            cfg = name_to_cfg.get(name)
+            spec = _spec_for(name)
+            cfg = spec[0] if spec else None
             if cfg is None:
                 row = (name, 0.0, float("inf"), "UNKNOWN_CONFIG")
             elif ctx_dead:
diff --git a/benchmark/gemm/frost/benchmark_moe_block_scale_matmul_swiglu.py b/benchmark/gemm/frost/benchmark_moe_block_scale_matmul_swiglu.py
index 4ca67b70b..d51b61424 100644
--- a/benchmark/gemm/frost/benchmark_moe_block_scale_matmul_swiglu.py
+++ b/benchmark/gemm/frost/benchmark_moe_block_scale_matmul_swiglu.py
@@ -9,6 +9,7 @@
 from __future__ import annotations
 
 import argparse
+import re
 import sys
 from typing import Callable
 
@@ -19,6 +20,7 @@
 from types import SimpleNamespace
 
 from cudnn.gemm.frost.compiler import jit_from_cudnn_graph
+from cudnn.gemm.frost.tile_config import by_name as _by_name
 from cudnn.gemm.frost.graph_analyzer import analyze
 from cudnn.gemm.frost.kernel_registry import candidates as _registry_candidates
 
@@ -244,6 +246,27 @@ def _build_spec_map():
 
 _SPEC_MAP = _build_spec_map()
 
+_LABEL_RE = re.compile(r"^(CONFIG_sm\d+_\d+x\d+x\d+_\d+x\d+x\d+_cluster\d+x\d+)_([12])ctamma(_static)?$")
+
+
+def _spec_for(name):
+    """(geometry cfg, cta_group, scheduler) for a --configs label, or None.
+
+    The sweep set comes from the registry funnel over CATALOG; a label naming a
+    geometry outside it (e.g. a num_mma_m > 1 tile, which `by_name` synthesizes) is
+    still runnable, so parse it rather than reporting it unsweepable."""
+    spec = _SPEC_MAP.get(name)
+    if spec is not None:
+        return spec
+    m = _LABEL_RE.match(name)
+    if m is None:
+        return None
+    try:
+        cfg = _by_name(m.group(1))
+    except (KeyError, NotImplementedError):
+        return None
+    return cfg, int(m.group(2)), "static" if m.group(3) else "clc"
+
 
 def main() -> int:
     p = argparse.ArgumentParser(description=__doc__, formatter_class=argparse.RawDescriptionHelpFormatter)
@@ -298,10 +321,11 @@ def main() -> int:
 
     best = None
     for label in config_names:
-        if label not in _SPEC_MAP:
+        spec = _spec_for(label)
+        if spec is None:
             print(f"  {label:66s} UNKNOWN (not a sweepable MoE block-scale swiglu strategy)")
             continue
-        cfg, cta_group, sched = _SPEC_MAP[label]
+        cfg, cta_group, sched = spec
         try:
             g, h = _graph_swiglu(S, N, K, E, combo)
             plan = _build_plan(g, cfg, cta_group, sched)
diff --git a/benchmark/gemm/frost/benchmark_moe_grouped_matmul.py b/benchmark/gemm/frost/benchmark_moe_grouped_matmul.py
index 921b8a0d3..f527c11c5 100644
--- a/benchmark/gemm/frost/benchmark_moe_grouped_matmul.py
+++ b/benchmark/gemm/frost/benchmark_moe_grouped_matmul.py
@@ -25,6 +25,7 @@
 import torch
 
 from cudnn.gemm.frost.compiler import jit_from_cudnn_graph
+from cudnn.gemm.frost.tile_config import by_name as _by_name
 from cudnn.gemm.frost.fusion_ir import (
     FusionChain as _FC,
     MatmulSpec as _MS,
@@ -43,7 +44,7 @@ def _vp_moe(handles, token, weight, fto, output):
 
 def _build_plan(g, cfg, name):
     """JIT-compile the recorded graph with a forced tile config."""
-    return jit_from_cudnn_graph(g, config=cfg, cta_group=_SPEC_MAP[name][1], scheduler=_SPEC_MAP[name][2])
+    return jit_from_cudnn_graph(g, config=cfg, cta_group=_spec_for(name)[1], scheduler=_spec_for(name)[2])
 
 
 def _build_spec_map():
@@ -72,6 +73,27 @@ def _build_spec_map():
 
 _SPEC_MAP = _build_spec_map()
 
+_LABEL_RE = re.compile(r"^(CONFIG_sm\d+_\d+x\d+x\d+_\d+x\d+x\d+_cluster\d+x\d+)_([12])ctamma(_static)?$")
+
+
+def _spec_for(name):
+    """(geometry cfg, cta_group, scheduler) for a --configs label, or None.
+
+    The sweep set comes from the registry funnel over CATALOG; a label naming a
+    geometry outside it (e.g. a num_mma_m > 1 tile, which `by_name` synthesizes) is
+    still runnable, so parse it rather than reporting UNKNOWN_CONFIG."""
+    spec = _SPEC_MAP.get(name)
+    if spec is not None:
+        return spec
+    m = _LABEL_RE.match(name)
+    if m is None:
+        return None
+    try:
+        cfg = _by_name(m.group(1))
+    except (KeyError, NotImplementedError):
+        return None
+    return cfg, int(m.group(2)), "static" if m.group(3) else "clc"
+
 
 # ---------------------------------------------------------------------------
 # Graph + data setup
@@ -377,10 +399,10 @@ def _nsys_worker(shape, configs, warmup, iters, nbuf) -> None:
     torch.cuda.synchronize()
 
     # 2. each MoE config.
-    name_to_cfg = {lbl: sp[0] for lbl, sp in _SPEC_MAP.items()}
     config_names = configs or list(_SPEC_MAP)
     for name in config_names:
-        cfg = name_to_cfg.get(name)
+        spec = _spec_for(name)
+        cfg = spec[0] if spec else None
         if cfg is None:
             continue
         try:
@@ -449,7 +471,6 @@ def main() -> int:
 
     flops = 2 * S * N * K
     config_names = [c.strip() for c in args.configs.split(",")] if args.configs else list(_SPEC_MAP)
-    name_to_cfg = {lbl: sp[0] for lbl, sp in _SPEC_MAP.items()}
 
     print(f"\n=== moe_grouped_matmul G={G} M={M} N={N} K={K}  " f"(S={S} tokens, ~{flops / 1e9:.1f} GFLOP) — BF16 ===")
 
@@ -482,11 +503,12 @@ def _fmt_row(name, tflops, ms, note, ref_tflops) -> str:
             cublas_tflops, cublas_ms = float("nan"), float("nan")
             print("  cuBLAS kernel: not detected in nsys output")
         for name in config_names:
-            cfg = name_to_cfg.get(name)
+            spec = _spec_for(name)
+            cfg = spec[0] if spec else None
             if cfg is None:
                 rows.append((name, 0.0, float("inf"), "UNKNOWN_CONFIG"))
                 continue
-            tok = _kernel_match_token(cfg, _SPEC_MAP[name][1])
+            tok = _kernel_match_token(cfg, _spec_for(name)[1])
             matches = [(k, v) for k, v in kern_times.items() if tok in k]
             if not matches:
                 rows.append((name, 0.0, float("inf"), "NO_KERNEL_IN_NSYS"))
@@ -519,7 +541,8 @@ def _fmt_row(name, tflops, ms, note, ref_tflops) -> str:
 
         ctx_dead = False
         for name in config_names:
-            cfg = name_to_cfg.get(name)
+            spec = _spec_for(name)
+            cfg = spec[0] if spec else None
             if cfg is None:
                 row = (name, 0.0, float("inf"), "UNKNOWN_CONFIG")
             elif ctx_dead:
diff --git a/benchmark/gemm/frost/benchmark_moe_grouped_matmul_models.py b/benchmark/gemm/frost/benchmark_moe_grouped_matmul_models.py
index 66cd5cf49..c9278dbc4 100644
--- a/benchmark/gemm/frost/benchmark_moe_grouped_matmul_models.py
+++ b/benchmark/gemm/frost/benchmark_moe_grouped_matmul_models.py
@@ -30,6 +30,7 @@
 from __future__ import annotations
 
 import argparse
+import re
 import sys
 from types import SimpleNamespace
 from typing import Callable
@@ -39,6 +40,7 @@
 import torch
 
 from cudnn.gemm.frost.compiler import jit_from_cudnn_graph
+from cudnn.gemm.frost.tile_config import by_name as _by_name
 from cudnn.gemm.frost.graph_analyzer import analyze
 from cudnn.gemm.frost.kernel_registry import candidates as _registry_candidates
 
@@ -399,6 +401,28 @@ def _build_spec_map(variant: str, dtype: str) -> dict[str, tuple]:
     return m
 
 
+_LABEL_RE = re.compile(r"^(CONFIG_sm\d+_\d+x\d+x\d+_\d+x\d+x\d+_cluster\d+x\d+)_([12])ctamma(_static)?$")
+
+
+def _spec_for(label: str, spec_map: dict):
+    """(geometry cfg, cta_group, scheduler) for a --configs label, or None.
+
+    Takes the sweep map because this one is built per (variant, dtype). A label
+    naming a geometry outside CATALOG (e.g. a num_mma_m > 1 tile, which `by_name`
+    synthesizes) is still runnable, so parse it rather than calling it unsweepable."""
+    spec = spec_map.get(label)
+    if spec is not None:
+        return spec
+    m = _LABEL_RE.match(label)
+    if m is None:
+        return None
+    try:
+        cfg = _by_name(m.group(1))
+    except (KeyError, NotImplementedError):
+        return None
+    return cfg, int(m.group(2)), "static" if m.group(3) else "clc"
+
+
 # Main
 
 
@@ -438,10 +462,11 @@ def _run_model(key: str, spec: dict, args) -> tuple | None:
 
     best = None
     for label in labels:
-        if label not in spec_map:
+        spec = _spec_for(label, spec_map)
+        if spec is None:
             print(f"  {label:64s} UNKNOWN (not a sweepable MoE dual-GEMM strategy)", flush=True)
             continue
-        cfg, cta_group, sched = spec_map[label]
+        cfg, cta_group, sched = spec
         try:
             g, h = _graph(S, N, K, E, variant, args.dtype, offsets)
             plan = jit_from_cudnn_graph(g, config=cfg, cta_group=cta_group, scheduler=sched)
diff --git a/benchmark/gemm/frost/benchmark_moe_grouped_matmul_swiglu.py b/benchmark/gemm/frost/benchmark_moe_grouped_matmul_swiglu.py
index 9418e43cc..a073cba99 100644
--- a/benchmark/gemm/frost/benchmark_moe_grouped_matmul_swiglu.py
+++ b/benchmark/gemm/frost/benchmark_moe_grouped_matmul_swiglu.py
@@ -11,6 +11,7 @@
 from __future__ import annotations
 
 import argparse
+import re
 import sys
 from typing import Callable
 
@@ -21,6 +22,7 @@
 from types import SimpleNamespace
 
 from cudnn.gemm.frost.compiler import jit_from_cudnn_graph
+from cudnn.gemm.frost.tile_config import by_name as _by_name
 from cudnn.gemm.frost.graph_analyzer import analyze
 from cudnn.gemm.frost.kernel_registry import candidates as _registry_candidates
 
@@ -197,6 +199,27 @@ def _build_spec_map():
 
 _SPEC_MAP = _build_spec_map()
 
+_LABEL_RE = re.compile(r"^(CONFIG_sm\d+_\d+x\d+x\d+_\d+x\d+x\d+_cluster\d+x\d+)_([12])ctamma(_static)?$")
+
+
+def _spec_for(name):
+    """(geometry cfg, cta_group, scheduler) for a --configs label, or None.
+
+    The sweep set comes from the registry funnel over CATALOG; a label naming a
+    geometry outside it (e.g. a num_mma_m > 1 tile, which `by_name` synthesizes) is
+    still runnable, so parse it rather than reporting it unsweepable."""
+    spec = _SPEC_MAP.get(name)
+    if spec is not None:
+        return spec
+    m = _LABEL_RE.match(name)
+    if m is None:
+        return None
+    try:
+        cfg = _by_name(m.group(1))
+    except (KeyError, NotImplementedError):
+        return None
+    return cfg, int(m.group(2)), "static" if m.group(3) else "clc"
+
 
 # Main
 
@@ -256,10 +279,11 @@ def main() -> int:
 
     best = None
     for label in config_names:
-        if label not in _SPEC_MAP:
+        spec = _spec_for(label)
+        if spec is None:
             print(f"  {label:64s} UNKNOWN (not a sweepable MoE swiglu strategy)")
             continue
-        cfg, cta_group, sched = _SPEC_MAP[label]
+        cfg, cta_group, sched = spec
         try:
             g, h = _graph_swiglu(S, N, K, E)
             plan = _build_plan(g, cfg, cta_group, sched)
diff --git a/include/cudnn_frontend_utils.h b/include/cudnn_frontend_utils.h
index 2ef2116ae..ebaa44610 100644
--- a/include/cudnn_frontend_utils.h
+++ b/include/cudnn_frontend_utils.h
@@ -702,6 +702,7 @@ enum class DataType_t {
     FP8_E5M2,
     FAST_FLOAT_FOR_FP8,
     FP8_E8M0,
+    FP8_E5M3,
     FP4_E2M1,
     INT4,
     COMPLEX_FP32,
@@ -728,6 +729,7 @@ NLOHMANN_JSON_SERIALIZE_ENUM(DataType_t,
                                  {DataType_t::FP8_E5M2, "FP8_E5M2"},
                                  {DataType_t::FAST_FLOAT_FOR_FP8, "FAST_FLOAT_FOR_FP8"},
                                  {DataType_t::FP8_E8M0, "FP8_E8M0"},
+                                 {DataType_t::FP8_E5M3, "FP8_E5M3"},
                                  {DataType_t::FP4_E2M1, "FP4_E2M1"},
                                  {DataType_t::INT4, "INT4"},
                                  {DataType_t::COMPLEX_FP32, "COMPLEX_FP32"},
@@ -1138,6 +1140,14 @@ convert_to_cudnn_type(cudnn_frontend::DataType_t const mode, cudnnDataType_t& cu
             return cudnnStatus_t::CUDNN_STATUS_SUCCESS;
 #else
             return cudnnStatus_t::CUDNN_STATUS_INVALID_VALUE;
+#endif
+        case DataType_t::FP8_E5M3:
+#if (CUDNN_VERSION >= 92600)
+            NV_CUDNN_FE_DYNAMIC_CHECK_CUDNN_BACKEND_VERSION(92600, cudnnStatus_t::CUDNN_STATUS_INVALID_VALUE);
+            cudnn_mode = CUDNN_DATA_FP8_E5M3;
+            return cudnnStatus_t::CUDNN_STATUS_SUCCESS;
+#else
+            return cudnnStatus_t::CUDNN_STATUS_INVALID_VALUE;
 #endif
         case DataType_t::FP4_E2M1:
 #if (CUDNN_VERSION >= 90700)
@@ -2349,6 +2359,10 @@ convert_from_cudnn_type(cudnnDataType_t const cudnn_mode) {
         case CUDNN_DATA_FP8_E8M0:
             return DataType_t::FP8_E8M0;
 #endif
+#if (CUDNN_VERSION >= 92600)
+        case CUDNN_DATA_FP8_E5M3:
+            return DataType_t::FP8_E5M3;
+#endif
 #if (CUDNN_VERSION >= 90700)
         case CUDNN_DATA_FP4_E2M1:
             return DataType_t::FP4_E2M1;
@@ -2410,6 +2424,9 @@ get_element_size_in_bits(cudnn_frontend::DataType_t datatype) {
 #endif
 #if (CUDNN_VERSION >= 90700)
         case DataType_t::FP8_E8M0:
+#endif
+#if (CUDNN_VERSION >= 92600)
+        case DataType_t::FP8_E5M3:
 #endif
             return 8;
             break;
diff --git a/python/cudnn/frost/device.py b/python/cudnn/frost/device.py
index c61771ddc..c66f60177 100644
--- a/python/cudnn/frost/device.py
+++ b/python/cudnn/frost/device.py
@@ -119,6 +119,22 @@ def shared_memory_per_block_optin(device: int) -> int:
     return int(_ck(*drv.cuDeviceGetAttribute(drv.CUdevice_attribute.CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK_OPTIN, handle)))
 
 
+# CU_DEVICE_ATTRIBUTE_MAX_OVERSIZED_SHARED_MEMORY_PER_BLOCK. Named in CUDA 13.4's
+# cuda.h; cuda-python's CUdevice_attribute does not carry it yet, so ask by ordinal.
+_ATTR_MAX_OVERSIZED_SHARED_MEMORY_PER_BLOCK = 150
+
+
+@functools.lru_cache(maxsize=None)
+def oversized_shared_memory_per_block(device: int) -> int:
+    """Per-CTA SMEM ceiling in the *oversized* carveout (327 KiB vs the 227 KiB
+    opt-in limit on SM 10.7), which the part gives by shrinking L1 to 8 kB — free for
+    a TMA-fed GEMM. 0 when the driver has no such mode."""
+    drv = _driver()
+    handle = _device_handle(device)
+    err, value = drv.cuDeviceGetAttribute(_ATTR_MAX_OVERSIZED_SHARED_MEMORY_PER_BLOCK, handle)
+    return int(value) if int(err) == 0 else 0
+
+
 @functools.lru_cache(maxsize=None)
 def l2_cache_bytes(device: int) -> int:
     drv = _driver()
diff --git a/python/cudnn/gemm/frost/compiler.py b/python/cudnn/gemm/frost/compiler.py
index 49e720979..4481268b8 100644
--- a/python/cudnn/gemm/frost/compiler.py
+++ b/python/cudnn/gemm/frost/compiler.py
@@ -51,6 +51,7 @@ def _as_custream(stream):
     _output_align_reqs,
     _pow2_floor,
     allowed_store_vsize,
+    dtype_arch_reject,
     tensor_alignment,
 )
 from .epilogue_codegen import EpilogueSnippets, generate
@@ -560,6 +561,8 @@ def _smem_desc_params(
     lines.append(f"gemm_b_idx = {tuple(b for _, b in chain.gemm_operands)}")
     total_tmem = _tmem_cols_for_arch()
     lines.append(f"num_tmem_alloc_cols = {total_tmem}")
+    lines.append(f"tmem_alloc_exclusive = {total_tmem > _MAX_NON_EXCLUSIVE_TMEM_COLS}")
+    lines.append(f"b_collector_ok = {_b_collector_supported()}")
     # TMEM accumulator budget. One acc stage holds, per GEMM, one region of
     # `num_mma_m` MMA-M blocks each `_epi_tile_cols` columns wide (the N-direction
     # MMAs subdivide that width, they do not add to it). `total_tmem == 0` means
@@ -667,35 +670,108 @@ def _smem_desc_params(
     return "\n".join(lines)
 
 
-def _quant_device_imports(chain: FusionChain) -> list[str]:
-    """fp32 -> ue8m0 scale byte via the sm100 cvt unit (round-up, satfinite),
-    emitted into the generated kernel so it stays self-contained. Semantics
-    match the 2^ceil(log2(x)) scale reference; x == 0 gives byte 0 (2^-127).
-    The DSL's own .to(Float8E8M0FNU) lowers to a ~9-instruction emulation."""
-    if not any(q.scale_dtype == "fp8_e8m0" for q in chain.quants):
-        return []
+def _cvt_f32_to_fp8_scale_bits(fn_name: str, dsl_dtype: str, rnd: str) -> list[str]:
+    """A ``fp32 -> <8-bit scale> byte`` helper, emitted into the generated kernel
+    so it stays self-contained. The x2 destination gets the value in both lanes;
+    only the low byte is read back."""
     return [
-        "from cutlass.cutlass_dsl import T as _frost_T",
-        "from cutlass._mlir.dialects import llvm as _frost_llvm",
         "",
         "",
-        "def _frost_cvt_f32_to_e8m0_bits(x):",
+        f"def {fn_name}(x):",
         "    src = cutlass.Float32(x).ir_value()",
-        '    asm = "{\\n  .reg .b16 lo;\\n  cvt.rp.satfinite.ue8m0x2.f32 lo, 0f00000000, $1;\\n  cvt.u32.u16 $0, lo;\\n}"',
-        "    byte = _frost_llvm.inline_asm(",
-        "        _frost_T.i32(),",
-        "        [src],",
-        "        asm,",
-        '        "=r,f",',
-        "        has_side_effects=False,",
-        "        is_align_stack=False,",
-        "        asm_dialect=_frost_llvm.AsmDialect.AD_ATT,",
+        "    pair = _frost_vector.from_elements(_frost_ir.VectorType.get([2], cutlass.Float32.mlir_type), [src, src])",
+        "    lo = _frost_vector.extract(pair, dynamic_position=[], static_position=[0])",
+        "    hi = _frost_vector.extract(pair, dynamic_position=[], static_position=[1])",
+        "    packed = _frost_nvvm.convert_f32x2_to_f8x2(",
+        "        _frost_ir.VectorType.get([2], cutlass.Int8.mlir_type),",
+        "        hi,",
+        "        lo,",
+        f"        _frost_ir.TypeAttr.get(cutlass.{dsl_dtype}.mlir_type),",
+        f"        rnd=_frost_nvvm.FPRoundingMode.{rnd},",
+        "        sat=_frost_nvvm.SaturationMode.SATFINITE,",
         "    )",
-        "    return cutlass.Int32(byte)",
+        "    byte = _frost_llvm.zext(_frost_T.i32(), _frost_llvm.bitcast(cutlass.Int16.mlir_type, packed))",
+        "    return cutlass.Int32(byte) & 0xFF",
     ]
 
 
-_TMEM_COLS_BY_ARCH: tuple[tuple[tuple[int, int], int], ...] = (((100, 120), 512),)
+def _cvt_e5m3_bits_to_f32() -> list[str]:
+    """The inverse of :func:`_cvt_f32_to_fp8_scale_bits` for ue5m3. It goes
+    through **bf16**, not fp16: bf16 carries 8 exponent bits, so it holds the
+    whole E5M3 range including bytes 248..254 (up to 114688), which are finite
+    because the format is canonical-NaN-only. fp16 would turn those into inf."""
+    return [
+        "",
+        "",
+        "def _frost_e5m3_bits_to_f32(b):",
+        "    byte = _frost_llvm.trunc(_frost_T.i8(), cutlass.Int32(b).ir_value(), _frost_llvm.IntegerOverflowFlags.none)",
+        "    pair = _frost_vector.from_elements(_frost_ir.VectorType.get([2], cutlass.Int8.mlir_type), [byte, byte])",
+        "    widened = _frost_nvvm.convert_f8x2_to_bf16x2(",
+        "        _frost_ir.VectorType.get([2], cutlass.BFloat16.mlir_type),",
+        "        pair,",
+        "        _frost_ir.TypeAttr.get(cutlass.FloatNV8E5M3FNU.mlir_type),",
+        "    )",
+        "    return cutlass.Float32(cutlass.BFloat16(_frost_vector.extract(widened, dynamic_position=[], static_position=[0])))",
+    ]
+
+
+def _quant_device_imports(chain: FusionChain) -> list[str]:
+    """Device-side converters between fp32 and the two scale formats whose
+    user-level DSL cast is not usable here. These reach the same hardware cvt
+    unit through the typed NVVM ops the cast itself is built on, which is what
+    lets them ask for the two things the cast does not do:
+
+    * ``sat=SATFINITE`` on the narrowing. The plain cast overflows to byte 255
+      (NaN); a NaN scale poisons its whole block on dequantize. Measured on
+      sm_107, that saturation is the ONLY way the cast differs here.
+    * ``ue5m3 -> bf16`` on the widening. The cast widens through a type that
+      reads E == 31 as inf, but E5M3 is canonical-NaN-only, so bytes 248..253
+      are finite (up to 114688) and come back as inf.
+
+    Both round UP: a scale rounded DOWN makes ``amax / scale`` exceed the output
+    format's max, clamping the block's largest element.
+
+    ``ue8m0`` needs no widening helper — it is a bare exponent, so ``byte << 23``
+    IS the fp32. ``ue5m3``'s cvt exists ONLY on sm_107, see the arch gate in
+    :func:`_check_block_quant_supported`.
+
+    Both take ``x == 0`` to byte 0, which the readback turns back into 0.0."""
+    kinds = {q.scale_dtype for q in chain.quants}
+    lines: list[str] = []
+    if "fp8_e8m0" in kinds:
+        lines += _cvt_f32_to_fp8_scale_bits("_frost_cvt_f32_to_e8m0_bits", "Float8E8M0FNU", "RP")
+    if "fp8_e5m3" in kinds:
+        lines += _cvt_f32_to_fp8_scale_bits("_frost_cvt_f32_to_e5m3_bits", "FloatNV8E5M3FNU", "RP")
+        lines += _cvt_e5m3_bits_to_f32()
+    if not lines:
+        return []
+    return [
+        "from cutlass.cutlass_dsl import T as _frost_T",
+        "from cutlass._mlir import ir as _frost_ir",
+        "from cutlass._mlir.dialects import llvm as _frost_llvm, nvvm as _frost_nvvm, vector as _frost_vector",
+    ] + lines
+
+
+# TMEM columns the GPU has — a HARDWARE property, so every pipeline running on
+# a given arch gets the same budget
+_TMEM_COLS_BY_ARCH: tuple[tuple[tuple[int, int], int], ...] = (
+    ((100, 107), 512),
+    ((107, 110), 576),
+    ((110, 120), 512),
+)
+
+# Past this, tcgen05.alloc must ask for the exclusive mode (and the count stops
+# being a power of two, so it goes through as a register operand).
+_MAX_NON_EXCLUSIVE_TMEM_COLS = 512
+
+
+_B_COLLECTOR_ARCH_RANGES: tuple[tuple[int, int], ...] = ((107, 110),)
+
+
+def _b_collector_supported(arch: int | None = None) -> bool:
+    """Whether this GPU's MMA can hold B in a collector buffer across MMAs."""
+    a = _current_arch() if arch is None else arch
+    return a is not None and any(lo <= a < hi for lo, hi in _B_COLLECTOR_ARCH_RANGES)
 
 
 def _tmem_cols_for_arch(arch: int | None = None) -> int:
@@ -775,8 +851,9 @@ def _render_block_scale_tile_constants(
     assert bs is not None
     is_fp4 = bs.is_fp4
     is_sm103 = cfg.pipeline == "sm103"
+    is_sm107 = cfg.pipeline == "sm107"
     if is_sm103 and not is_fp4:
-        raise NotImplementedError("the sm103 block-scale pipeline is fp4-only (nvfp4/mxfp4); " f"{bs.combo} runs the sm100 templates")
+        raise NotImplementedError("the sm103 block-scale pipeline is fp4-only; " f"{bs.a_dtype} data with {bs.sf_dtype} scales runs the sm100 templates")
 
     # data_elem_bits / sA-sB bytes / B's TMA stride encoding all take one packed
     # width, read off A alone — a mixed-width combo would mis-size B, not fail.
@@ -1011,10 +1088,13 @@ def _align16(x: int) -> int:
     out_dt = chain.output_dtype
     vec_bytes_epi = _epi_vec_bytes(chain, cfg, cta_group)
 
-    # Instruction-descriptor operand dtype. fp4 MMA uses Tcgen05MxInstrDesc with
-    # the E5M2 piggy-back; fp8 uses the real fp8 dtype.
-    if is_fp4:
+    # Instruction-descriptor operand dtype. On sm100 the fp4 MMA rides
+    # Tcgen05MxInstrDesc with the E5M2 piggy-back; the K=64B sm107 fp4 MMA is an
+    # OMMA and takes the real fp4 dtype. fp8 always uses its real dtype.
+    if is_fp4 and not is_sm107:
         idesc_a = idesc_b = "cutlass.Float8E5M2"
+    elif is_fp4:
+        idesc_a = idesc_b = "cutlass.Float4E2M1FN"
     else:
         idesc_a = DTYPE_TO_CUTLASS[bs.a_dtype]
         idesc_b = DTYPE_TO_CUTLASS[bs.b_dtype]
@@ -1023,7 +1103,7 @@ def _align16(x: int) -> int:
     ab_tma_format = "_tma.TensorMapDataFormat.B4X16" if is_fp4 else "None"
 
     lines = [
-        f"# Block-scale config: {cfg.name} combo={bs.combo}",
+        f"# Block-scale config: {cfg.name} data={bs.a_dtype}x{bs.b_dtype} sf={bs.sf_dtype} block={bs.block_size}",
         f"cta_tile_m = {cta_m}",
         f"cta_tile_n = {cta_n}",
         f"cta_tile_k_elems = {cta_k_elems}",
@@ -1144,6 +1224,8 @@ def _align16(x: int) -> int:
         f"sfa_col_base = {sfa_col_base}",
         f"sfb_col_base = {sfb_col_base}",
         f"num_tmem_alloc_cols = {num_tmem_alloc_cols}",
+        f"tmem_alloc_exclusive = {num_tmem_alloc_cols > _MAX_NON_EXCLUSIVE_TMEM_COLS}",
+        f"b_collector_ok = {_b_collector_supported()}",
         f"sfa_smem_bytes = {sfa_smem_bytes}",
         f"sfb_smem_bytes = {sfb_smem_bytes}",
         f"sf_tma_box_k = {sf_tma_box_k}",
@@ -1176,6 +1258,18 @@ def _align16(x: int) -> int:
             f"sfa_mma_col_off_by_j = {tuple(spi * j // 4 * 4 * nb_m for j in range(num_kblocks))}",
             f"sfb_mma_col_off_by_j = {tuple(spi * j // 4 * 4 * nb_n for j in range(num_kblocks))}",
         ]
+    if is_sm107:
+        # SM 10.7 block-scale MMA: K = 64 bytes per instruction (2x sm100), so
+        # one MMA consumes sf_scales_per_inst scales — 8 for nvfp4, which spans
+        # word_atoms = 2 of the 4-scale 128x4 utccp atoms. fp4 is an OMMA
+        # (K-mode 2 = 128 fp4 elements); mxfp8 stays on the MX descriptor
+        # (K-mode 1 = 64 fp8 elements).
+        lines += [
+            "",
+            f"# sm107 K=64B block-scale MMA: {num_kblocks} MMAs per K-tile",
+            f"idesc_is_omma = {is_fp4}",
+            f"mma_k_dim_mode = {2 if is_fp4 else 1}",
+        ]
     # MoE grouped block-scale: grouped persistent scheduler launches a FIXED
     # cluster count (≈ NUM_SMS / cluster_size); host grid and stride share it.
     # first_token_offset dtype (int32/int64) drives the compile() fake.
@@ -2585,6 +2679,9 @@ def probe_supported(
     Block-scale / MoE gate inside their ``_jit_*`` compile paths; here a
     successful analysis is treated as eligible (full validation at compile)."""
     chain, _binding = analyze_with_binding(graph)
+    _dtype_reason = dtype_arch_reject(chain, _current_arch())
+    if _dtype_reason is not None:
+        raise NotImplementedError(_dtype_reason)
     _check_executable(chain)
     if chain.has_moe or chain.has_block_scale:
         return  # specialized paths validate at compile
@@ -2601,7 +2698,7 @@ def probe_supported(
     _check_supported(chain, config)
     from .kernel_registry import select_template as _sel_tmpl
 
-    _arch_reason = _sel_tmpl(chain, config, cta_group, scheduler).arch_active_reject()
+    _arch_reason = _sel_tmpl(chain, config, cta_group, scheduler).active_reject(config)
     if _arch_reason is not None:
         raise NotImplementedError(_arch_reason)
     _check_dtype_config_compat(chain, config, cta_group)
@@ -2628,6 +2725,9 @@ def jit_from_cudnn_graph(
     ``force_stg_epi=True`` skips the TMA-store path even when its gate accepts.
     """
     chain, binding = analyze_with_binding(graph)
+    _dtype_reason = dtype_arch_reject(chain, _current_arch())
+    if _dtype_reason is not None:
+        raise NotImplementedError(_dtype_reason)
     _check_cta_group_geometry(config, cta_group)
     _check_mma_n_dim(chain, config, cta_group)
     # MoE grouped block-scale = both matches at once (dequant + moe_grouped);
@@ -2659,7 +2759,7 @@ def jit_from_cudnn_graph(
     _check_supported(chain, config)
     from .kernel_registry import select_template as _sel_tmpl
 
-    _arch_reason = _sel_tmpl(chain, config, cta_group, scheduler).arch_active_reject()
+    _arch_reason = _sel_tmpl(chain, config, cta_group, scheduler).active_reject(config)
     if _arch_reason is not None:
         raise NotImplementedError(_arch_reason)
     _check_dtype_config_compat(chain, config, cta_group)
@@ -3074,7 +3174,7 @@ def _jit_moe(
         )
     from .kernel_registry import select_template as _sel_tmpl
 
-    _arch_reason = _sel_tmpl(chain, config, cta_group, scheduler).arch_active_reject()
+    _arch_reason = _sel_tmpl(chain, config, cta_group, scheduler).active_reject(config)
     if _arch_reason is not None:
         raise NotImplementedError(_arch_reason)
     _check_dtype_config_compat(chain, config, cta_group)
@@ -3130,14 +3230,17 @@ def _jit_block_scale(
     # both-sided requirement (single-sided matches no case).
     _check_block_scale_supported(chain, config.pipeline)
     _check_input_alignment(chain)
+    # The sm107 templates carry the sm100 epilogue verbatim, so reductions and
+    # the quant epilogue ride it unchanged; sm103's own pipeline does not.
+    _epi_pipelines = ("sm100", "sm107")
     if chain.reductions:
-        if config.pipeline != "sm100":
-            raise NotImplementedError("block-scale reduction is supported only on sm100 templates")
+        if config.pipeline not in _epi_pipelines:
+            raise NotImplementedError(f"block-scale reduction is not supported on {config.pipeline} templates")
         for red in chain.reductions:
             if red.compute_dtype != "fp32" or red.dtype != "fp32":
                 raise NotImplementedError("block-scale reduction supports only fp32 compute/output")
-    if chain.quants and config.pipeline != "sm100":
-        raise NotImplementedError("block-scale quant epilogue is supported only on sm100 templates " "(not yet validated on sm103)")
+    if chain.quants and config.pipeline not in _epi_pipelines:
+        raise NotImplementedError(f"block-scale quant epilogue is not supported on {config.pipeline} " "templates (not yet validated on sm103)")
     # Per-template active-GPU SM gate (no-op when no GPU is visible).
     from .kernel_registry import select_template
 
@@ -3146,7 +3249,7 @@ def _jit_block_scale(
         raise NotImplementedError(
             f"block-scale multi-GEMM ({chain.num_gemms} GEMMs) is not supported by " f"{_tmpl.file} (cta_group={cta_group}, scheduler={scheduler!r})."
         )
-    _arch_reason = _tmpl.arch_active_reject()
+    _arch_reason = _tmpl.active_reject(config)
     if _arch_reason is not None:
         raise NotImplementedError(_arch_reason)
     _compute_output_vec_bytes(chain)  # eager: rejects bad output alignment
@@ -3486,7 +3589,7 @@ def _jit_moe_block_scale(
     _check_input_alignment(chain)
     # Per-template active-GPU SM gate (no-op when no GPU is visible).
     _tmpl = select_template(chain, config, cta_group, scheduler)
-    _arch_reason = _tmpl.arch_active_reject()
+    _arch_reason = _tmpl.active_reject(config)
     if _arch_reason is not None:
         raise NotImplementedError(_arch_reason)
     _compute_output_vec_bytes(chain)
diff --git a/python/cudnn/gemm/frost/dtypes.py b/python/cudnn/gemm/frost/dtypes.py
index 460792e9f..e06d7ac8c 100644
--- a/python/cudnn/gemm/frost/dtypes.py
+++ b/python/cudnn/gemm/frost/dtypes.py
@@ -25,6 +25,11 @@
     "fp8_e4m3": "cutlass.Float8E4M3FN",
     "fp8_e5m2": "cutlass.Float8E5M2",
     "fp8_e8m0": "cutlass.Float8E8M0FNU",
+    # E5M3 only ever appears as a scale factor, which the kernel takes as a base
+    # pointer — the format lives in the MMA descriptor, not the tensor type. So
+    # it rides an opaque byte (cutlass.FloatNV8E5M3FNU exists but TVM-FFI cannot
+    # marshal it, and torch has no E5M3 dtype for the runtime buffer either).
+    "fp8_e5m3": "cutlass.Uint8",
     "fp4_e2m1": "cutlass.Float4E2M1FNx2",
     "uint8": "cutlass.Uint8",
     "int32": "cutlass.Int32",
@@ -40,6 +45,7 @@
     "fp8_e4m3": 1,
     "fp8_e5m2": 1,
     "fp8_e8m0": 1,
+    "fp8_e5m3": 1,
     "fp4_e2m1": 1,
     "uint8": 1,
     "int32": 4,
@@ -50,6 +56,29 @@
 # is stored packed 2/byte, so DTYPE_BYTES reads 1 and cannot tell fp4 from fp8.
 DTYPE_BITS: dict[Dtype, int] = {**{dt: nbytes * 8 for dt, nbytes in DTYPE_BYTES.items()}, "fp4_e2m1": 4}
 
+DTYPE_GPU_ARCH_RANGES: dict[Dtype, tuple[tuple[int, int], ...]] = {
+    "fp8_e5m3": ((107, 110),),
+}
+
+
+def _fmt_ranges(ranges: tuple[tuple[int, int], ...]) -> str:
+    return " or ".join(f"{lo} <= SM < {hi}" for lo, hi in ranges)
+
+
+def dtype_arch_reject(chain: FusionChain, arch: "int | None") -> "str | None":
+    """Why the active GPU cannot run this chain's dtypes, or ``None``.
+
+    ``arch`` is ``None`` when no GPU is visible (render-only / CI), which skips
+    the check the same way the other arch gates do."""
+    if arch is None:
+        return None
+    for dtype in sorted(chain.dtypes_used()):
+        ranges = DTYPE_GPU_ARCH_RANGES.get(dtype)
+        if ranges is not None and not any(lo <= arch < hi for lo, hi in ranges):
+            return f"dtype {dtype!r} exists only on {_fmt_ranges(ranges)}, but the active GPU is sm_{arch}"
+    return None
+
+
 # input dtype -> tcgen05 MMA kind.
 DTYPE_TO_MMA_KIND: dict[Dtype, str] = {
     "bf16": "nvvm.Tcgen05MMAKind.F16",
@@ -68,6 +97,7 @@
     cudnn.data_type.FP8_E4M3: "fp8_e4m3",
     cudnn.data_type.FP8_E5M2: "fp8_e5m2",
     cudnn.data_type.FP8_E8M0: "fp8_e8m0",
+    cudnn.data_type.FP8_E5M3: "fp8_e5m3",
     cudnn.data_type.FP4_E2M1: "fp4_e2m1",
     cudnn.data_type.UINT8: "uint8",
     cudnn.data_type.INT32: "int32",
diff --git a/python/cudnn/gemm/frost/epilogue_codegen.py b/python/cudnn/gemm/frost/epilogue_codegen.py
index ed3c0b6b6..a5580e430 100644
--- a/python/cudnn/gemm/frost/epilogue_codegen.py
+++ b/python/cudnn/gemm/frost/epilogue_codegen.py
@@ -832,6 +832,43 @@ def _quant_output_min(dtype: Dtype) -> str:
     raise ValueError(f"block quantize output dtype {dtype!r} is not supported by codegen")
 
 
+def _scale_store_dtype(scale_dtype: Dtype) -> str:
+    """The DSL type a quantized scale is STORED as — one source of truth for the
+    scale tap's element type, its zero-init, and the value written. E5M3 has no
+    DSL float type and a raw_ptr store to a Uint8 tensor is rejected, so it rides
+    the Int8 byte carrier (the same one packed FP4 data uses)."""
+    return "cutlass.Int8" if scale_dtype == "fp8_e5m3" else DTYPE_TO_CUTLASS[scale_dtype]
+
+
+def _emit_scale_quantize(p: str, sfx: str, src: str, scale_var: str, back_var: str, quant: BlockQuantizeSpec) -> list[str]:
+    """Quantize one fp32 scale to ``quant.scale_dtype`` and read the STORED
+    value back as fp32 — the data is divided by what was actually written, not
+    by the pre-rounding scale, so a dequantize reproduces it exactly.
+
+    E4M3 round-trips through the DSL ``.to()``. The other two reach the cvt unit
+    through the helpers :func:`compiler._quant_device_imports` emits (which
+    documents why their ``.to()`` is not usable), and read the byte back as
+    ``byte << 23`` for ue8m0 — a bare exponent, so that IS the fp32, and byte 0
+    is 0.0 — or through the paired widening helper for ue5m3."""
+    scale_dtype = _scale_store_dtype(quant.scale_dtype)
+    if quant.scale_dtype == "fp8_e8m0":
+        return [
+            f"{p}_qb{sfx} = _frost_cvt_f32_to_e8m0_bits({src})",
+            f"{scale_var} = (({p}_qb{sfx}).to(cutlass.Int8)).bitcast({scale_dtype})",
+            f"{back_var} = ({p}_qb{sfx} << 23).bitcast(cutlass.Float32)",
+        ]
+    if quant.scale_dtype == "fp8_e5m3":
+        return [
+            f"{p}_qb{sfx} = _frost_cvt_f32_to_e5m3_bits({src})",
+            f"{scale_var} = ({p}_qb{sfx}).to({scale_dtype})",
+            f"{back_var} = _frost_e5m3_bits_to_f32({p}_qb{sfx})",
+        ]
+    return [
+        f"{scale_var} = ({src}).to({scale_dtype})",
+        f"{back_var} = ({scale_var}).to(cutlass.Float32)",
+    ]
+
+
 def _emit_block_quant_col(
     quant: BlockQuantizeSpec,
     quant_idx: int,
@@ -849,7 +886,7 @@ def _emit_block_quant_col(
     stores the scale byte(s) of column(s) ``col_j + k*G + l % G``. The
     compiler gates the row guards to be reduction-uniform."""
     p = f"_q{quant_idx}"
-    scale_dtype = DTYPE_TO_CUTLASS[quant.scale_dtype]
+    scale_dtype = _scale_store_dtype(quant.scale_dtype)
     G = 16 if quant.block_size == 16 else 32
     if vsize % G != 0:
         raise NotImplementedError(
@@ -889,24 +926,7 @@ def _emit_block_quant_col(
                 lines.append(f'{p}_a{i} = cute.arch.warp_redux_sync({p}_src[{i}], "fmax", abs=True)')
         for i in cols:
             lines.append(f"{p}_s{i} = {p}_a{i} * {p}_rl")
-            if quant.scale_dtype == "fp8_e8m0":
-                # HW cvt (rp/satfinite) instead of the ~9-instruction emulated
-                # .to(E8M0); the widened value only feeds rcp, so byte<<23
-                # (0.0 for byte 0 -> rcp inf -> FLT_MAX clamp) is equivalent.
-                lines.extend(
-                    [
-                        f"{p}_qb{i} = _frost_cvt_f32_to_e8m0_bits({p}_s{i})",
-                        f"{p}_q{i} = (({p}_qb{i}).to(cutlass.Int8)).bitcast({scale_dtype})",
-                        f"{p}_u{i} = ({p}_qb{i} << 23).bitcast(cutlass.Float32)",
-                    ]
-                )
-            else:
-                lines.extend(
-                    [
-                        f"{p}_q{i} = ({p}_s{i}).to({scale_dtype})",
-                        f"{p}_u{i} = ({p}_q{i}).to(cutlass.Float32)",
-                    ]
-                )
+            lines.extend(_emit_scale_quantize(p, str(i), f"{p}_s{i}", f"{p}_q{i}", f"{p}_u{i}", quant))
             lines.extend(
                 [
                     f"{p}_i{i} = cute.math.min(cute.arch.rcp_approx({p}_u{i}), cutlass.Float32(3.402823466e38))",
@@ -993,7 +1013,6 @@ def _emit_block_quant(
     if vsize % bs != 0:
         raise NotImplementedError(f"row block-quantize: store vector {vsize} must be a multiple of block_size {bs}")
     n_sub = vsize // bs
-    scale_dtype = DTYPE_TO_CUTLASS[quant.scale_dtype]
     lines: list[str] = [
         f"{p}_src = ({source_var}).to(cutlass.Float32)",
         f"{p}_abs = cute.math.abs({p}_src)",
@@ -1006,21 +1025,7 @@ def _emit_block_quant(
         for e in range(1, bs):
             lines.append(f"{p}_amax{k} = cute.math.max({p}_amax{k}, {p}_abs[{base + e}])")
         lines.append(f"{p}_sf{k} = {p}_amax{k} * {p}_rl")
-        if quant.scale_dtype == "fp8_e8m0":
-            lines.extend(
-                [
-                    f"{p}_qb{k} = _frost_cvt_f32_to_e8m0_bits({p}_sf{k})",
-                    f"{p}_scale{k} = (({p}_qb{k}).to(cutlass.Int8)).bitcast({scale_dtype})",
-                    f"{p}_up{k} = ({p}_qb{k} << 23).bitcast(cutlass.Float32)",
-                ]
-            )
-        else:
-            lines.extend(
-                [
-                    f"{p}_scale{k} = ({p}_sf{k}).to({scale_dtype})",
-                    f"{p}_up{k} = ({p}_scale{k}).to(cutlass.Float32)",
-                ]
-            )
+        lines.extend(_emit_scale_quantize(p, str(k), f"{p}_sf{k}", f"{p}_scale{k}", f"{p}_up{k}", quant))
         lines.append(f"{p}_inv{k} = cute.math.min(cute.arch.rcp_approx({p}_up{k}), cutlass.Float32(3.402823466e38))")
         for e in range(bs):
             lines.append(f"{p}_out[{base + e}] = {p}_src[{base + e}] * {p}_inv{k}")
@@ -1322,7 +1327,11 @@ def _scale_tap_idx(qi: int) -> int:
     # a false claim for reduction / quant-scale / M-major taps.
     _out_reqs = _output_align_reqs(chain, use_tma_store, vec_bytes=vec_bytes_epi)
     for i, tap in enumerate(taps):
-        _fake_dt = "cutlass.Int8" if (not tap.is_reduction and not tap.is_quant_scale and tap.dtype == "fp4_e2m1") else DTYPE_TO_CUTLASS[tap.dtype]
+        # Byte-carrier taps: packed FP4 data, and an E5M3 scale (the DSL has no
+        # E5M3 float type, and a raw_ptr store to a Uint8 tensor is rejected —
+        # Int8 is the 8-bit carrier this package already uses for FP4).
+        _fp4_data_tap = not tap.is_reduction and not tap.is_quant_scale and tap.dtype == "fp4_e2m1"
+        _fake_dt = "cutlass.Int8" if _fp4_data_tap else _scale_store_dtype(tap.dtype) if tap.is_quant_scale else DTYPE_TO_CUTLASS[tap.dtype]
         # Every tap is consumed as a raw pointer (gC_tap_i_ptr) plus explicit
         # out_/red_/quant_scale_stride_* scalars, so the only genuine layout
         # contract is stride_n == 1 on an N-major DENSE tap (_dense_store_offset).
diff --git a/python/cudnn/gemm/frost/fusion_ir.py b/python/cudnn/gemm/frost/fusion_ir.py
index 772f6e1e6..b4e01d9db 100644
--- a/python/cudnn/gemm/frost/fusion_ir.py
+++ b/python/cudnn/gemm/frost/fusion_ir.py
@@ -15,6 +15,7 @@
 
 from __future__ import annotations
 
+import dataclasses
 from dataclasses import dataclass, field
 from typing import Literal
 
@@ -27,6 +28,7 @@
     "fp8_e4m3",
     "fp8_e5m2",
     "fp8_e8m0",
+    "fp8_e5m3",
     "fp4_e2m1",
     "int8",
     "uint8",
@@ -40,6 +42,7 @@
     "fp8_e4m3",
     "fp8_e5m2",
     "fp8_e8m0",
+    "fp8_e5m3",
     "fp4_e2m1",
     "int8",
     "uint8",
@@ -50,10 +53,6 @@
 BMajor = Literal["k", "n"]
 OutMajor = Literal["n", "m"]
 
-# Block-scaled-matmul scale-factor dtypes (SFA/SFB): nvfp4 = FP4 data + E4M3
-# scale block16; mxfp4 = FP4 + E8M0 block32; mxfp8 = FP8 + E8M0 block32.
-BLOCK_SCALE_SF_DTYPES: tuple[Dtype, ...] = ("fp8_e4m3", "fp8_e8m0")
-
 # Aux broadcast onto the (M, N) tile: scalar / per_row (len M) / per_col
 # (len N) / per_elem (full M×N).
 BroadcastMode = Literal["scalar", "per_row", "per_col", "per_elem"]
@@ -257,8 +256,8 @@ def __post_init__(self) -> None:
             raise NotImplementedError(f"block quantize supports the N axis (-1/2) or the M axis (1) in " f"cudnn.gemm.frost; got axis={self.axis}")
         if self.transpose and self.axis != 1:
             raise ValueError("block quantize transpose=True requires the M axis (axis=1)")
-        if self.scale_dtype not in ("fp8_e8m0", "fp8_e4m3"):
-            raise ValueError(f"block quantize scale dtype {self.scale_dtype!r} is not supported; " "expected fp8_e8m0 or fp8_e4m3")
+        if self.scale_dtype not in ("fp8_e8m0", "fp8_e4m3", "fp8_e5m3"):
+            raise ValueError(f"block quantize scale dtype {self.scale_dtype!r} is not supported; " "expected fp8_e8m0, fp8_e4m3 or fp8_e5m3")
         if self.scale_reorder not in (None, "F8_128x4"):
             raise ValueError(f"block quantize scale reordering {self.scale_reorder!r} is not supported; " "expected None or F8_128x4")
         if self.compute_dtype != "fp32":
@@ -440,8 +439,10 @@ class BlockScaleSpec:
     one block-scale matmul (three shapes: dequant(A)@B, A@dequant(B),
     dequant(A)@dequant(B); ``sfa``/``sfb`` present only for the scaled side(s)).
     No dtype/block/arch rules here — runnability is decided at compile time.
-    Currently runs (both sides): nvfp4 (fp4+e4m3, block16), mxfp4 (fp4+e8m0,
-    block32), mxfp8 (fp8+e8m0, block32).
+    Currently runs (both sides): fp4 with any of e4m3 / e8m0 / e5m3 scales at
+    either K-block (16 or 32) — the two axes are orthogonal, so nvfp4 and mxfp4
+    are just the two best-known corners — plus fp8 (e4m3/e5m2) with e8m0 scales
+    at block 32. E5M3 scales additionally require SM 10.7+.
 
     SF tensors are runtime-positional (not ``TensorRef``s), fully described here
     by per-side scalars; their logical dims derive from M/N/K/block_size. Passed
@@ -497,13 +498,6 @@ def block_size(self) -> int:
     def sf_dtype(self) -> "Dtype | None":
         return self.sf_dtype_a if self.sf_dtype_a is not None else self.sf_dtype_b
 
-    @property
-    def combo(self) -> str:
-        """One of 'nvfp4', 'mxfp4', 'mxfp8' — the dtype/block family."""
-        if self.a_dtype == "fp4_e2m1":
-            return "nvfp4" if self.sf_dtype == "fp8_e4m3" else "mxfp4"
-        return "mxfp8"
-
     @property
     def is_fp4(self) -> bool:
         return self.a_dtype == "fp4_e2m1"
@@ -521,9 +515,16 @@ def scale_vec_size(self) -> str:
 
     @property
     def sf_scale_format(self) -> int:
-        """``Tcgen05MxInstrDesc`` ``scale_format`` field: 0 for E4M3 scale
-        (nvfp4), 1 for E8M0 scale (mx)."""
-        return 0 if self.sf_dtype == "fp8_e4m3" else 1
+        """``Tcgen05MxInstrDesc`` ``scale_format`` field (idesc bits 24-23);
+        E5M3 (2) needs SM 10.7+ silicon. Unlike :attr:`combo` this one is BAKED
+        INTO THE KERNEL, so an unregistered SF dtype must not fall back to some
+        other format's bits — that would miscompute silently. Declines instead
+        (``NotImplementedError`` is what the engine router treats as "does not
+        serve this graph"; a ``KeyError`` would escape as an engine bug)."""
+        fmt = {"fp8_e4m3": 0, "fp8_e8m0": 1, "fp8_e5m3": 2}.get(self.sf_dtype)
+        if fmt is None:
+            raise NotImplementedError(f"no MMA scale_format encoding for block-scale SF dtype {self.sf_dtype!r}")
+        return fmt
 
 
 @dataclass(frozen=True)
@@ -555,6 +556,23 @@ def __post_init__(self) -> None:
             raise ValueError(f"first_token_offset dtype must be int32 or int64; " f"got {self.offset_dtype!r}")
 
 
+def _walk_dtype_fields(obj: object, found: "set[Dtype]", *, in_dtype_field: bool = False) -> None:
+    """Recursive half of :meth:`FusionChain.dtypes_used`. A field counts when its
+    declared type mentions ``Dtype``; containers inherit that from their field."""
+    if isinstance(obj, str):
+        if in_dtype_field:
+            found.add(obj)
+    elif dataclasses.is_dataclass(obj) and not isinstance(obj, type):
+        for f in dataclasses.fields(obj):
+            _walk_dtype_fields(getattr(obj, f.name), found, in_dtype_field="Dtype" in str(f.type))
+    elif isinstance(obj, dict):
+        for item in obj.values():
+            _walk_dtype_fields(item, found, in_dtype_field=in_dtype_field)
+    elif isinstance(obj, (list, tuple, set, frozenset)):
+        for item in obj:
+            _walk_dtype_fields(item, found, in_dtype_field=in_dtype_field)
+
+
 @dataclass
 class FusionChain:
     """Full description of a matmul + pointwise-epilogue fusion (linear chain
@@ -785,6 +803,14 @@ def taps(self) -> list["ChainOutput"]:
         output exists — reductions/scale still ride the tap plumbing)."""
         return self.outputs[1:] if self.output_specs else list(self.outputs)
 
+    def dtypes_used(self) -> frozenset[Dtype]:
+        """Every dtype this chain names, anywhere. Found by walking the IR's
+        ``Dtype``-annotated fields rather than listing the dtype-bearing specs,
+        so a gate built on this cannot be bypassed by a field added later."""
+        found: set[Dtype] = set()
+        _walk_dtype_fields(self, found)
+        return frozenset(found)
+
     def summary(self) -> str:
         """One-line human-readable summary for logs / error messages."""
         m = self.matmul
diff --git a/python/cudnn/gemm/frost/graph_analyzer.py b/python/cudnn/gemm/frost/graph_analyzer.py
index af37ccbe4..8ce2c5c22 100644
--- a/python/cudnn/gemm/frost/graph_analyzer.py
+++ b/python/cudnn/gemm/frost/graph_analyzer.py
@@ -563,7 +563,8 @@ def build_gemm_plan(graph: cudnn.pygraph):
     if not _graph_has_gemm(graph):
         raise ValueError("cudnn.gemm.frost: graph has no matmul / moe_grouped_matmul node; nothing to compile")
     from .compiler import jit_from_cudnn_graph
-    from .tile_config import select_config
+    from .kernel_registry import preferred_pipeline
+    from .tile_config import as_pipeline, select_config
 
     chain = analyze(graph)
     tile_m = chain.matmul.M
@@ -587,6 +588,7 @@ def build_gemm_plan(graph: cudnn.pygraph):
         b_elem_bytes=DTYPE_BYTES[chain.matmul.b_dtype],
         supports_static=supports_static,
     )
+    config = as_pipeline(config, preferred_pipeline(chain))
     return jit_from_cudnn_graph(graph, config=config, cta_group=cta_group, scheduler=scheduler)
 
 
diff --git a/python/cudnn/gemm/frost/kernel_registry.py b/python/cudnn/gemm/frost/kernel_registry.py
index e8e7b34fd..82df5f1a2 100644
--- a/python/cudnn/gemm/frost/kernel_registry.py
+++ b/python/cudnn/gemm/frost/kernel_registry.py
@@ -41,14 +41,11 @@ def _pipeline_from_file(template_file: str) -> str:
 
 
 # Active-GPU SM ranges per template pipeline — half-open [lo, hi) segments,
-# SM = major*10 + minor. A family may support several DISJOINT segments (e.g.
-# a future family running on sm200-210 plus sm250-270), so this is a tuple of
-# segments, never a single lo/hi pair.
+# SM = major*10 + minor. A family may support several DISJOINT segments
 PIPELINE_ARCH_RANGES: dict[str, tuple[tuple[int, int], ...]] = {
-    # sm100 templates use only family-portable Blackwell instructions.
     "sm100": ((100, 120),),
-    # The fp4 K=48B is an arch-exact ("a"-level) feature of SM 10.3.
-    "sm103": ((103, 104),),
+    "sm103": ((103, 110),),
+    "sm107": ((107, 110),),
 }
 
 # Pointwise ops a mainloop-fusion template can transform in SMEM.
@@ -141,16 +138,18 @@ def _bs_key(a: str, sfa: str, b: str, sfb: str, kblk: int) -> tuple:
     )
 
 
-# Supported block-scale (data, SF dtype, K-block) cases — shared by the plain
-# block-scale matmul and the block-scaled MoE grouped matmul.
 _BLOCK_SCALE_CASES = frozenset(
     {
-        _bs_key("fp4_e2m1", "fp8_e4m3", "fp4_e2m1", "fp8_e4m3", 16),  # nvfp4
-        _bs_key("fp4_e2m1", "fp8_e8m0", "fp4_e2m1", "fp8_e8m0", 32),  # mxfp4
-        _bs_key("fp8_e4m3", "fp8_e8m0", "fp8_e4m3", "fp8_e8m0", 32),  # mxfp8 e4m3×e4m3
-        _bs_key("fp8_e4m3", "fp8_e8m0", "fp8_e5m2", "fp8_e8m0", 32),  # mxfp8 e4m3×e5m2
-        _bs_key("fp8_e5m2", "fp8_e8m0", "fp8_e4m3", "fp8_e8m0", 32),  # mxfp8 e5m2×e4m3
-        _bs_key("fp8_e5m2", "fp8_e8m0", "fp8_e5m2", "fp8_e8m0", 32),  # mxfp8 e5m2×e5m2
+        _bs_key("fp4_e2m1", "fp8_e4m3", "fp4_e2m1", "fp8_e4m3", 16),
+        _bs_key("fp4_e2m1", "fp8_e4m3", "fp4_e2m1", "fp8_e4m3", 32),
+        _bs_key("fp4_e2m1", "fp8_e8m0", "fp4_e2m1", "fp8_e8m0", 16),
+        _bs_key("fp4_e2m1", "fp8_e8m0", "fp4_e2m1", "fp8_e8m0", 32),
+        _bs_key("fp4_e2m1", "fp8_e5m3", "fp4_e2m1", "fp8_e5m3", 16),
+        _bs_key("fp4_e2m1", "fp8_e5m3", "fp4_e2m1", "fp8_e5m3", 32),
+        _bs_key("fp8_e4m3", "fp8_e8m0", "fp8_e4m3", "fp8_e8m0", 32),
+        _bs_key("fp8_e4m3", "fp8_e8m0", "fp8_e5m2", "fp8_e8m0", 32),
+        _bs_key("fp8_e5m2", "fp8_e8m0", "fp8_e4m3", "fp8_e8m0", 32),
+        _bs_key("fp8_e5m2", "fp8_e8m0", "fp8_e5m2", "fp8_e8m0", 32),
     }
 )
 
@@ -193,11 +192,18 @@ def _bs_key(a: str, sfa: str, b: str, sfb: str, kblk: int) -> tuple:
     "sm103": {
         GraphType.BLOCK_SCALE_MATMUL: frozenset(
             {
-                _bs_key("fp4_e2m1", "fp8_e4m3", "fp4_e2m1", "fp8_e4m3", 16),  # nvfp4
-                _bs_key("fp4_e2m1", "fp8_e8m0", "fp4_e2m1", "fp8_e8m0", 32),  # mxfp4
+                _bs_key("fp4_e2m1", "fp8_e4m3", "fp4_e2m1", "fp8_e4m3", 16),
+                _bs_key("fp4_e2m1", "fp8_e4m3", "fp4_e2m1", "fp8_e4m3", 32),
+                _bs_key("fp4_e2m1", "fp8_e8m0", "fp4_e2m1", "fp8_e8m0", 16),
+                _bs_key("fp4_e2m1", "fp8_e8m0", "fp4_e2m1", "fp8_e8m0", 32),
+                _bs_key("fp4_e2m1", "fp8_e5m3", "fp4_e2m1", "fp8_e5m3", 16),
+                _bs_key("fp4_e2m1", "fp8_e5m3", "fp4_e2m1", "fp8_e5m3", 32),
             }
         ),
     },
+    "sm107": {
+        GraphType.BLOCK_SCALE_MATMUL: _BLOCK_SCALE_CASES,
+    },
 }
 
 # The ONE home for checks that need template SM family × mma dtype × ACTUAL
@@ -206,10 +212,16 @@ def _bs_key(a: str, sfa: str, b: str, sfb: str, kblk: int) -> tuple:
 # combos never appear here (stage-0 family gate + the existence sets decide).
 # Values are half-open [lo, hi) segments, same shape as PIPELINE_ARCH_RANGES.
 MMA_GPU_ARCH_SPECIAL_CASES: dict[tuple[str, tuple], tuple[tuple[int, int], ...]] = {
-    # int8 UTCIMMA (tcgen05 kind::i8) exists only on SM 100 and SM 110 — the
-    # other Blackwell family members dropped it (e.g. sm103 reworked the
-    # tensor core for fp4 throughput).
     ("sm100", ("int8", "int8", "int32")): ((100, 101), (110, 111)),
+    ("sm100", _bs_key("fp4_e2m1", "fp8_e5m3", "fp4_e2m1", "fp8_e5m3", 16)): ((107, 110),),
+    ("sm100", _bs_key("fp4_e2m1", "fp8_e5m3", "fp4_e2m1", "fp8_e5m3", 32)): ((107, 110),),
+    ("sm100", _bs_key("fp4_e2m1", "fp8_e4m3", "fp4_e2m1", "fp8_e4m3", 32)): ((107, 110),),
+    ("sm103", _bs_key("fp4_e2m1", "fp8_e5m3", "fp4_e2m1", "fp8_e5m3", 16)): ((107, 110),),
+    ("sm103", _bs_key("fp4_e2m1", "fp8_e5m3", "fp4_e2m1", "fp8_e5m3", 32)): ((107, 110),),
+    ("sm103", _bs_key("fp4_e2m1", "fp8_e4m3", "fp4_e2m1", "fp8_e4m3", 32)): ((107, 110),),
+    ("sm107", _bs_key("fp4_e2m1", "fp8_e5m3", "fp4_e2m1", "fp8_e5m3", 16)): ((107, 110),),
+    ("sm107", _bs_key("fp4_e2m1", "fp8_e5m3", "fp4_e2m1", "fp8_e5m3", 32)): ((107, 110),),
+    ("sm107", _bs_key("fp4_e2m1", "fp8_e4m3", "fp4_e2m1", "fp8_e4m3", 32)): ((107, 110),),
 }
 
 
@@ -259,6 +271,8 @@ class KernelTemplate:
     mainloop: bool  # mainloop-fusion variant (transform A/B before MMA)
     # Multi-GEMM support (templates without it reject multi-GEMM chains).
     supports_multi_gemm: bool = False
+    # A CTA tile spanning several MMA instructions along M (num_mma_m > 1).
+    supports_multi_mma_m: bool = True
 
     @property
     def block_scale(self) -> bool:
@@ -313,6 +327,9 @@ def _config_reject(self, chain: FusionChain, config: TileConfig) -> str | None:
             C._check_mma_n_dim(chain, config, self.cta_group)
         except NotImplementedError as e:
             return str(e)
+        reason = self.multi_mma_m_reject(config)
+        if reason is not None:
+            return reason
         try:
             if self.block_scale:
                 from .tile_config import validate_block_scale_config
@@ -369,6 +386,23 @@ def accepts(self, chain: FusionChain, config: TileConfig) -> str | None:
             or self._other_reject(chain, config)
         )
 
+    def multi_mma_m_reject(self, config: TileConfig) -> str | None:
+        """``None`` unless the config splits the CTA tile across several MMA
+        instructions along M and this template has not been adapted to it."""
+        if config.num_mma_m > 1 and not self.supports_multi_mma_m:
+            return (
+                f"{self.file} does not support a CTA tile spanning several MMA "
+                f"instructions along M (num_mma_m={config.num_mma_m}); use a "
+                f"cta_tile_m of {config.mma_inst_m}"
+            )
+        return None
+
+    def active_reject(self, config: TileConfig) -> str | None:
+        """The gates a JIT path applies once it has picked this template: the
+        active GPU's SM range, then capabilities a pure-geometry config can ask
+        for that this template does not implement."""
+        return self.arch_active_reject() or self.multi_mma_m_reject(config)
+
     def candidate_configs(self, chain: FusionChain) -> tuple[TileConfig, ...]:
         """Catalog geometries this template accepts for ``chain`` — by
         predicate filter, never hand-maintained."""
@@ -388,7 +422,7 @@ def _extra_reject(self, chain: FusionChain, config: TileConfig) -> str | None:
         return None
 
 
-# Registry — one entry per template file (15 today). A geometry config expands
+# Registry — one entry per template file (20 today). A geometry config expands
 # across these via `candidates`. cta_group / static_sched / mainloop live HERE.
 
 
@@ -400,6 +434,7 @@ def _mm(
     mainloop: bool = False,
     graph_type: GraphType = GraphType.MATMUL,
     supports_multi_gemm: bool = False,
+    supports_multi_mma_m: bool = True,
 ) -> KernelTemplate:
     pipeline = _pipeline_from_file(file)
     if pipeline not in PIPELINE_ARCH_RANGES:
@@ -415,6 +450,7 @@ def _mm(
         graph_type=graph_type,
         mainloop=mainloop,
         supports_multi_gemm=supports_multi_gemm,
+        supports_multi_mma_m=supports_multi_mma_m,
     )
 
 
@@ -465,17 +501,37 @@ def _mm(
     ),
     # sm103 block-scaled matmul: fp4-only (nvfp4/mxfp4), K=48B UTCOMMA
     # (K-tile 384 B, 8 MMAs over 3× 128-B chunks via circular SMEM descs).
+    # num_mma_m > 1 is NOT adapted here: the chunk pipeline miscomputes (A reads
+    # unwritten SMEM in K, period 192 B) and the ab_stages budget under-counts,
+    # so cta_tile_m=256 also overruns the SMEM cap. Both are silent-wrong /
+    # launch-fail, hence the gate. See CLAUDE.md for what was ruled out.
     _mm(
         "sm103_block_scale_matmul_1ctamma.py",
         cta_group=1,
         static=False,
         graph_type=GraphType.BLOCK_SCALE_MATMUL,
+        supports_multi_mma_m=False,
     ),
     _mm(
         "sm103_block_scale_matmul_2ctamma.py",
         cta_group=2,
         static=False,
         graph_type=GraphType.BLOCK_SCALE_MATMUL,
+        supports_multi_mma_m=False,
+    ),
+    _mm(
+        "sm107_block_scale_matmul_1ctamma.py",
+        cta_group=1,
+        static=False,
+        graph_type=GraphType.BLOCK_SCALE_MATMUL,
+        supports_multi_gemm=True,
+    ),
+    _mm(
+        "sm107_block_scale_matmul_2ctamma.py",
+        cta_group=2,
+        static=False,
+        graph_type=GraphType.BLOCK_SCALE_MATMUL,
+        supports_multi_gemm=True,
     ),
     # mainloop-fusion matmul (CLC only — no static / block-scale variant yet)
     _mm("sm100_matmul_mainloop_1ctamma.py", cta_group=1, static=False, mainloop=True),
@@ -511,9 +567,41 @@ def _mm(
         graph_type=GraphType.MOE_BLOCK_SCALE,
         supports_multi_gemm=True,
     ),
+    _mm(
+        "sm107_moe_grouped_block_scale_matmul_fwd_1ctamma.py",
+        cta_group=1,
+        static=False,
+        graph_type=GraphType.MOE_BLOCK_SCALE,
+        supports_multi_gemm=True,
+    ),
+    _mm(
+        "sm107_moe_grouped_block_scale_matmul_fwd_2ctamma.py",
+        cta_group=2,
+        static=False,
+        graph_type=GraphType.MOE_BLOCK_SCALE,
+        supports_multi_gemm=True,
+    ),
 )
 
 
+# Pipeline families the AUTO path (``tile_config.select_config``) may build
+# with, best first. sm103 is deliberately absent: its 384-byte K-tile is outside
+# select_config's geometry ladder, so it stays an explicit-config pipeline.
+_AUTO_PIPELINE_ORDER: tuple[str, ...] = ("sm107", "sm100")
+
+
+def preferred_pipeline(chain: FusionChain) -> str:
+    """Pipeline family the auto path should build ``chain`` with: the first
+    :data:`_AUTO_PIPELINE_ORDER` entry that has a template for this graph type
+    and whose SM range covers the active GPU. A graph type the newer family
+    does not implement (plain matmul, MoE) falls through to sm100 by itself."""
+    gt = classify_graph_type(chain)
+    for pipeline in _AUTO_PIPELINE_ORDER:
+        if any(t.pipeline == pipeline and t.graph_type is gt and t.arch_active_reject() is None for t in TEMPLATES):
+            return pipeline
+    return _AUTO_PIPELINE_ORDER[-1]
+
+
 def select_template(
     chain: FusionChain,
     config: TileConfig,
diff --git a/python/cudnn/gemm/frost/kernel_templates/_tile_helpers.py b/python/cudnn/gemm/frost/kernel_templates/_tile_helpers.py
new file mode 100644
index 000000000..bffc67a9f
--- /dev/null
+++ b/python/cudnn/gemm/frost/kernel_templates/_tile_helpers.py
@@ -0,0 +1,131 @@
+# Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: MIT
+
+"""Tile-level helpers shared by the rendered kernel templates.
+
+A template is RENDERED (its `@@INJECT_*@@` blocks become module-level
+constants) and then exec'd from the kernel cache under a synthetic module
+name, so it cannot use relative imports and this module is never rendered.
+Everything here therefore takes what it needs as ARGUMENTS -- a helper that
+reads an injected constant (`num_mma_m`, `tile_swizzle_n`, `ab_dtype`, ...)
+has to stay in the template, or be re-signed to receive it.
+"""
+
+from __future__ import annotations
+
+import cutlass
+import cutlass.cute as cute
+import cutlass.experimental.primitives as nvvm
+
+
+@cute.jit
+def l2_swizzle_tile(raw_m, raw_n, nt_m, nt_n, swizzle_w):
+    """N-direction super-block rasterization of the (m, n) cgrp-tile coord, for
+    L2 reuse. ``swizzle_w == 1`` falls out of the math as the identity mapping.
+    """
+    t = raw_n * nt_m + raw_m
+    blk = nt_m * swizzle_w
+    sb = t // blk
+    off = t - sb * blk
+    base_n = sb * swizzle_w
+    cur_S = cutlass.min(cutlass.Int32(swizzle_w), nt_n - base_n)
+    log_m = off // cur_S
+    log_n = base_n + off - log_m * cur_S
+    return log_m, log_n
+
+
+def epi_subtile_spans(cols):
+    """Power-of-two column spans the epilogue drains a tile in (host-side)."""
+    spans = []
+    off = 0
+    while off < cols:
+        w = 32
+        while w > cols - off:
+            w //= 2
+        spans.append((off, w))
+        off += w
+    return spans
+
+
+TENSOR_MAP_QWORDS = 16
+
+
+def moe_swizzle_tile(t, nt_m, nt_n, swizzle_w):
+    """Group-local linear tile index -> (m, n) under an N-super-block walk.
+    ``swizzle_w == nt_n`` reproduces the plain n-fast split; ``1`` gives m-fast.
+    """
+    blk = cutlass.max(nt_m * swizzle_w, cutlass.Int32(1))
+    sb = t // blk
+    off = t - sb * blk
+    base_n = sb * swizzle_w
+    cur_S = cutlass.min(cutlass.Int32(swizzle_w), nt_n - base_n)
+    tile_m = off // cur_S
+    tile_n = base_n + off - tile_m * cur_S
+    return tile_m, tile_n
+
+
+@cute.jit
+def replace_tensormap_global_dim_1(desc_ptr, new_dim) -> None:
+    nvvm.tensormap_replace(
+        nvvm.TensormapField.GLOBAL_DIM,
+        desc_ptr,
+        new_value=cutlass.Int32(new_dim),
+        ord=1,
+    )
+
+
+@cute.jit
+def replace_tensormap_global_address(desc_ptr, new_address) -> None:
+    nvvm.tensormap_replace(
+        nvvm.TensormapField.GLOBAL_ADDRESS,
+        desc_ptr,
+        new_value=cutlass.Int64(new_address),
+    )
+
+
+@cute.jit
+def fence_tensormap_release() -> None:
+    nvvm.fence_proxy_release(
+        nvvm.MemScope.GPU,
+        from_proxy=nvvm.Proxy.GENERIC,
+        to_proxy=nvvm.Proxy.TENSORMAP,
+    )
+
+
+@cute.jit
+def fence_tensormap_acquire(desc_ptr) -> None:
+    nvvm.fence_proxy_acquire(
+        nvvm.MemScope.GPU,
+        desc_ptr,
+        TENSOR_MAP_QWORDS * 8,
+        from_proxy=nvvm.Proxy.GENERIC,
+        to_proxy=nvvm.Proxy.TENSORMAP,
+    )
+
+
+@cute.jit
+def moe_group_at(visit_idx, num_groups, num_experts):
+    """Visitation index -> routed group index.
+
+    ``num_groups == num_experts`` (or a non-multiple) walks groups in order. Batched MoE
+    (``num_groups == B * num_experts``) walks expert-major -- the B groups sharing expert
+    ``g % E`` become consecutive, so the expert weight is fetched once instead of B times.
+    """
+    per_expert = num_groups // cutlass.max(num_experts, cutlass.Int32(1))
+    group = visit_idx
+    if per_expert > 1 and per_expert * num_experts == num_groups:
+        group = (visit_idx % per_expert) * num_experts + (visit_idx // per_expert)
+    return group
+
+
+@cute.jit
+def copy_tensormap_to_workspace(src_desc_ptr, dst_i64_ptr) -> None:
+    """Copy the 128-byte A tensormap into ``dst_i64_ptr`` (seeds the SMEM copy).
+
+    The trip count is a compile-time constant, so this is a constexpr loop --
+    the templates had drifted into two spellings of the same fully-unrolled
+    copy (`range_constexpr` vs `range(..., unroll_full=True)`).
+    """
+    src_words = cute.make_ptr(cutlass.Int64, src_desc_ptr.toint(), mem_space=cute.AddressSpace.generic)
+    for i in cutlass.range_constexpr(TENSOR_MAP_QWORDS):
+        dst_i64_ptr.subview(i).store((src_words + i).load())
diff --git a/python/cudnn/gemm/frost/kernel_templates/sm100_block_scale_matmul_1ctamma.py b/python/cudnn/gemm/frost/kernel_templates/sm100_block_scale_matmul_1ctamma.py
index 6b0c190d1..bfbd1378c 100644
--- a/python/cudnn/gemm/frost/kernel_templates/sm100_block_scale_matmul_1ctamma.py
+++ b/python/cudnn/gemm/frost/kernel_templates/sm100_block_scale_matmul_1ctamma.py
@@ -23,6 +23,9 @@
 from typing import Callable
 
 import cutlass.experimental.primitives as nvvm
+from cudnn.gemm.frost.kernel_templates._tile_helpers import (
+    l2_swizzle_tile as _l2_swizzle_tile,
+)
 import cutlass.experimental.cuda.tensor_map as _tma
 import cutlass._mlir_helpers.vector as _cvec
 from cutlass import apply_swizzle as _apply_smem_swizzle
@@ -72,19 +75,18 @@ def _auto_swizzle_w(m, n, k, nt_n):
     return cutlass.Int32(w)
 
 
-def _l2_swizzle_tile(raw_m, raw_n, nt_m, nt_n, swizzle_w):
-    """N-direction super-block rasterization of the (m, n) cgrp-tile coord, for
-    L2 reuse. ``swizzle_w == 1`` falls out of the math as the identity mapping.
-    """
-    t = raw_n * nt_m + raw_m
-    blk = nt_m * swizzle_w
-    sb = t // blk
-    off = t - sb * blk
-    base_n = sb * swizzle_w
-    cur_S = cutlass.min(cutlass.Int32(swizzle_w), nt_n - base_n)
-    log_m = off // cur_S
-    log_n = base_n + off - log_m * cur_S
-    return log_m, log_n
+def _b_collector_op(mi):
+    """B is identical across the M sub-blocks (only A's address advances), so the
+    first MMA fills the B collector and the rest read it back instead of
+    re-fetching the same operand from SMEM. `.collector::b::*` is silicon-gated
+    (sm_107a only), hence `b_collector_ok`."""
+    if cutlass.const_expr(not b_collector_ok or num_mma_m == 1):
+        return None
+    if cutlass.const_expr(mi == 0):
+        return nvvm.Tcgen05MMACollectorOp.FILL
+    if cutlass.const_expr(mi == num_mma_m - 1):
+        return nvvm.Tcgen05MMACollectorOp.LASTUSE
+    return nvvm.Tcgen05MMACollectorOp.USE
 
 
 @cute.kernel
@@ -606,7 +608,12 @@ def _kernel(
 
     if warp_idx == mma_warp_id:
         nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
-        nvvm.tcgen05_alloc(tmem_ptr_i32, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_1)
+        nvvm.tcgen05_alloc(
+            tmem_ptr_i32,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_1,
+        )
         nvvm.bar_warp_sync(0xFFFFFFFF)
         nvvm.barrier_cta_arrive(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
         tmem_raw_addr = tmem_ptr_i32.load()
@@ -766,6 +773,7 @@ def _kernel(
                                         scale_a=sfa_dst_ptrs[_ai][mi],
                                         scale_b=sfb_scale_ptrs[_bj],
                                         scale_vec_size=scale_vec_size,
+                                        b_collector_op=_b_collector_op(mi),
                                     )
                         # Every accumulator sees scale_d=False on exactly the first
                         # k_block of the tile, so the flip stays outside mi.
@@ -823,7 +831,12 @@ def _kernel(
 
         nvvm.bar_warp_sync(0xFFFFFFFF)
         alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-        nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_1)
+        nvvm.tcgen05_dealloc(
+            alloc_ptr,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_1,
+        )
 
     if warp_idx < num_epilogue_warps:
         nvvm.setmaxregister(epi_reg_count, nvvm.SetMaxRegisterAction.INCREASE)
diff --git a/python/cudnn/gemm/frost/kernel_templates/sm100_block_scale_matmul_1ctamma_static.py b/python/cudnn/gemm/frost/kernel_templates/sm100_block_scale_matmul_1ctamma_static.py
index 149eb9c63..15a125691 100644
--- a/python/cudnn/gemm/frost/kernel_templates/sm100_block_scale_matmul_1ctamma_static.py
+++ b/python/cudnn/gemm/frost/kernel_templates/sm100_block_scale_matmul_1ctamma_static.py
@@ -25,6 +25,9 @@
 from typing import Callable
 
 import cutlass.experimental.primitives as nvvm
+from cudnn.gemm.frost.kernel_templates._tile_helpers import (
+    l2_swizzle_tile as _l2_swizzle_tile,
+)
 import cutlass.experimental.cuda.tensor_map as _tma
 import cutlass._mlir_helpers.vector as _cvec
 from cutlass import apply_swizzle as _apply_smem_swizzle
@@ -71,19 +74,18 @@ def _auto_swizzle_w(m, n, k, nt_n):
     return cutlass.Int32(w)
 
 
-def _l2_swizzle_tile(raw_m, raw_n, nt_m, nt_n, swizzle_w):
-    """N-direction super-block rasterization of the (m, n) cgrp-tile coord, for
-    L2 reuse. ``swizzle_w == 1`` falls out of the math as the identity mapping.
-    """
-    t = raw_n * nt_m + raw_m
-    blk = nt_m * swizzle_w
-    sb = t // blk
-    off = t - sb * blk
-    base_n = sb * swizzle_w
-    cur_S = cutlass.min(cutlass.Int32(swizzle_w), nt_n - base_n)
-    log_m = off // cur_S
-    log_n = base_n + off - log_m * cur_S
-    return log_m, log_n
+def _b_collector_op(mi):
+    """B is identical across the M sub-blocks (only A's address advances), so the
+    first MMA fills the B collector and the rest read it back instead of
+    re-fetching the same operand from SMEM. `.collector::b::*` is silicon-gated
+    (sm_107a only), hence `b_collector_ok`."""
+    if cutlass.const_expr(not b_collector_ok or num_mma_m == 1):
+        return None
+    if cutlass.const_expr(mi == 0):
+        return nvvm.Tcgen05MMACollectorOp.FILL
+    if cutlass.const_expr(mi == num_mma_m - 1):
+        return nvvm.Tcgen05MMACollectorOp.LASTUSE
+    return nvvm.Tcgen05MMACollectorOp.USE
 
 
 @cute.kernel
@@ -511,7 +513,12 @@ def _kernel(
 
     if warp_idx == mma_warp_id:
         nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
-        nvvm.tcgen05_alloc(tmem_ptr_i32, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_1)
+        nvvm.tcgen05_alloc(
+            tmem_ptr_i32,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_1,
+        )
         nvvm.bar_warp_sync(0xFFFFFFFF)
         nvvm.barrier_cta_arrive(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
         tmem_raw_addr = tmem_ptr_i32.load()
@@ -672,6 +679,7 @@ def _kernel(
                                         scale_a=sfa_dst_ptrs[_ai][mi],
                                         scale_b=sfb_scale_ptrs[_bj],
                                         scale_vec_size=scale_vec_size,
+                                        b_collector_op=_b_collector_op(mi),
                                     )
                         # Every accumulator sees scale_d=False on exactly the first
                         # k_block of the tile, so the flip stays outside mi.
@@ -713,7 +721,12 @@ def _kernel(
 
         nvvm.bar_warp_sync(0xFFFFFFFF)
         alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-        nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_1)
+        nvvm.tcgen05_dealloc(
+            alloc_ptr,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_1,
+        )
 
     if warp_idx < num_epilogue_warps:
         nvvm.setmaxregister(epi_reg_count, nvvm.SetMaxRegisterAction.INCREASE)
diff --git a/python/cudnn/gemm/frost/kernel_templates/sm100_block_scale_matmul_2ctamma.py b/python/cudnn/gemm/frost/kernel_templates/sm100_block_scale_matmul_2ctamma.py
index b0096696a..4314cb39b 100644
--- a/python/cudnn/gemm/frost/kernel_templates/sm100_block_scale_matmul_2ctamma.py
+++ b/python/cudnn/gemm/frost/kernel_templates/sm100_block_scale_matmul_2ctamma.py
@@ -22,6 +22,9 @@
 from typing import Callable
 
 import cutlass.experimental.primitives as nvvm
+from cudnn.gemm.frost.kernel_templates._tile_helpers import (
+    l2_swizzle_tile as _l2_swizzle_tile,
+)
 import cutlass.experimental.cuda.tensor_map as _tma
 import cutlass._mlir_helpers.vector as _cvec
 from cutlass import apply_swizzle as _apply_smem_swizzle
@@ -72,19 +75,18 @@ def _auto_swizzle_w(m, n, k, nt_n):
     return cutlass.Int32(w)
 
 
-def _l2_swizzle_tile(raw_m, raw_n, nt_m, nt_n, swizzle_w):
-    """N-direction super-block rasterization of the (m, n) cgrp-tile coord, for
-    L2 reuse. ``swizzle_w == 1`` falls out of the math as the identity mapping.
-    """
-    t = raw_n * nt_m + raw_m
-    blk = nt_m * swizzle_w
-    sb = t // blk
-    off = t - sb * blk
-    base_n = sb * swizzle_w
-    cur_S = cutlass.min(cutlass.Int32(swizzle_w), nt_n - base_n)
-    log_m = off // cur_S
-    log_n = base_n + off - log_m * cur_S
-    return log_m, log_n
+def _b_collector_op(mi):
+    """B is identical across the M sub-blocks (only A's address advances), so the
+    first MMA fills the B collector and the rest read it back instead of
+    re-fetching the same operand from SMEM. `.collector::b::*` is silicon-gated
+    (sm_107a only), hence `b_collector_ok`."""
+    if cutlass.const_expr(not b_collector_ok or num_mma_m == 1):
+        return None
+    if cutlass.const_expr(mi == 0):
+        return nvvm.Tcgen05MMACollectorOp.FILL
+    if cutlass.const_expr(mi == num_mma_m - 1):
+        return nvvm.Tcgen05MMACollectorOp.LASTUSE
+    return nvvm.Tcgen05MMACollectorOp.USE
 
 
 @cute.kernel
@@ -633,7 +635,12 @@ def _kernel(
     ab_empty_arrive_mask = cutlass.Int16(a_part | b_part)
     if warp_idx == mma_warp_id:
         nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
-        nvvm.tcgen05_alloc(tmem_ptr_i32, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+        nvvm.tcgen05_alloc(
+            tmem_ptr_i32,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_2,
+        )
         nvvm.bar_warp_sync(0xFFFFFFFF)
         nvvm.barrier_cta_arrive(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
         tmem_raw_addr = tmem_ptr_i32.load()
@@ -801,6 +808,7 @@ def _kernel(
                                             scale_a=sfa_dst_ptrs[_ai][mi],
                                             scale_b=sfb_scale_ptrs[_bj],
                                             scale_vec_size=scale_vec_size,
+                                            b_collector_op=_b_collector_op(mi),
                                         )
                             # Every accumulator sees scale_d=False on exactly the first
                             # k_block of the tile, so the flip stays outside mi.
@@ -866,7 +874,12 @@ def _kernel(
             if cutlass.const_expr(not use_acc_overlap):
                 nvvm.mbarrier_arrive(peer_mbar, scope=nvvm.MemScope.CLUSTER, relaxed=True)
             alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-            nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+            nvvm.tcgen05_dealloc(
+                alloc_ptr,
+                cutlass.Int32(num_tmem_alloc_cols),
+                is_exclusive=tmem_alloc_exclusive,
+                group=nvvm.CTAGroup.CTA_2,
+            )
         else:
             tile_iter = cutlass.Int32(0)
             is_valid = cutlass.Int32(1)
@@ -901,7 +914,12 @@ def _kernel(
             while not nvvm.mbarrier_try_wait_parity(tmem_dealloc_mbar_ptr, 0, time_limit=10_000_000):
                 pass
             alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-            nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+            nvvm.tcgen05_dealloc(
+                alloc_ptr,
+                cutlass.Int32(num_tmem_alloc_cols),
+                is_exclusive=tmem_alloc_exclusive,
+                group=nvvm.CTAGroup.CTA_2,
+            )
 
     if warp_idx < num_epilogue_warps:
         nvvm.setmaxregister(epi_reg_count, nvvm.SetMaxRegisterAction.INCREASE)
diff --git a/python/cudnn/gemm/frost/kernel_templates/sm100_block_scale_matmul_2ctamma_static.py b/python/cudnn/gemm/frost/kernel_templates/sm100_block_scale_matmul_2ctamma_static.py
index e666eaec2..5906c2117 100644
--- a/python/cudnn/gemm/frost/kernel_templates/sm100_block_scale_matmul_2ctamma_static.py
+++ b/python/cudnn/gemm/frost/kernel_templates/sm100_block_scale_matmul_2ctamma_static.py
@@ -23,6 +23,9 @@
 from typing import Callable
 
 import cutlass.experimental.primitives as nvvm
+from cudnn.gemm.frost.kernel_templates._tile_helpers import (
+    l2_swizzle_tile as _l2_swizzle_tile,
+)
 import cutlass.experimental.cuda.tensor_map as _tma
 import cutlass._mlir_helpers.vector as _cvec
 from cutlass import apply_swizzle as _apply_smem_swizzle
@@ -69,19 +72,18 @@ def _auto_swizzle_w(m, n, k, nt_n):
     return cutlass.Int32(w)
 
 
-def _l2_swizzle_tile(raw_m, raw_n, nt_m, nt_n, swizzle_w):
-    """N-direction super-block rasterization of the (m, n) cgrp-tile coord, for
-    L2 reuse. ``swizzle_w == 1`` falls out of the math as the identity mapping.
-    """
-    t = raw_n * nt_m + raw_m
-    blk = nt_m * swizzle_w
-    sb = t // blk
-    off = t - sb * blk
-    base_n = sb * swizzle_w
-    cur_S = cutlass.min(cutlass.Int32(swizzle_w), nt_n - base_n)
-    log_m = off // cur_S
-    log_n = base_n + off - log_m * cur_S
-    return log_m, log_n
+def _b_collector_op(mi):
+    """B is identical across the M sub-blocks (only A's address advances), so the
+    first MMA fills the B collector and the rest read it back instead of
+    re-fetching the same operand from SMEM. `.collector::b::*` is silicon-gated
+    (sm_107a only), hence `b_collector_ok`."""
+    if cutlass.const_expr(not b_collector_ok or num_mma_m == 1):
+        return None
+    if cutlass.const_expr(mi == 0):
+        return nvvm.Tcgen05MMACollectorOp.FILL
+    if cutlass.const_expr(mi == num_mma_m - 1):
+        return nvvm.Tcgen05MMACollectorOp.LASTUSE
+    return nvvm.Tcgen05MMACollectorOp.USE
 
 
 @cute.kernel
@@ -533,7 +535,12 @@ def _kernel(
     ab_empty_arrive_mask = cutlass.Int16(a_part | b_part)
     if warp_idx == mma_warp_id:
         nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
-        nvvm.tcgen05_alloc(tmem_ptr_i32, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+        nvvm.tcgen05_alloc(
+            tmem_ptr_i32,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_2,
+        )
         nvvm.bar_warp_sync(0xFFFFFFFF)
         nvvm.barrier_cta_arrive(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
         tmem_raw_addr = tmem_ptr_i32.load()
@@ -701,6 +708,7 @@ def _kernel(
                                             scale_a=sfa_dst_ptrs[_ai][mi],
                                             scale_b=sfb_scale_ptrs[_bj],
                                             scale_vec_size=scale_vec_size,
+                                            b_collector_op=_b_collector_op(mi),
                                         )
                             # Every accumulator sees scale_d=False on exactly the first
                             # k_block of the tile, so the flip stays outside mi.
@@ -751,7 +759,12 @@ def _kernel(
             if cutlass.const_expr(not use_acc_overlap):
                 nvvm.mbarrier_arrive(peer_mbar, scope=nvvm.MemScope.CLUSTER, relaxed=True)
             alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-            nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+            nvvm.tcgen05_dealloc(
+                alloc_ptr,
+                cutlass.Int32(num_tmem_alloc_cols),
+                is_exclusive=tmem_alloc_exclusive,
+                group=nvvm.CTAGroup.CTA_2,
+            )
         else:
             if cutlass.const_expr(USE_PDL):
                 if elect_one:
@@ -764,7 +777,12 @@ def _kernel(
             while not nvvm.mbarrier_try_wait_parity(tmem_dealloc_mbar_ptr, 0, time_limit=10_000_000):
                 pass
             alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-            nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+            nvvm.tcgen05_dealloc(
+                alloc_ptr,
+                cutlass.Int32(num_tmem_alloc_cols),
+                is_exclusive=tmem_alloc_exclusive,
+                group=nvvm.CTAGroup.CTA_2,
+            )
 
     if warp_idx < num_epilogue_warps:
         nvvm.setmaxregister(epi_reg_count, nvvm.SetMaxRegisterAction.INCREASE)
diff --git a/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_1ctamma.py b/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_1ctamma.py
index 2b865e706..2dc69605a 100644
--- a/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_1ctamma.py
+++ b/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_1ctamma.py
@@ -25,6 +25,10 @@
 from typing import Callable
 
 import cutlass.experimental.primitives as nvvm
+from cudnn.gemm.frost.kernel_templates._tile_helpers import (
+    epi_subtile_spans as _epi_subtile_spans,
+    l2_swizzle_tile as _l2_swizzle_tile,
+)
 import cutlass.experimental.cuda.tensor_map as _tma
 import cutlass._mlir_helpers.vector as _cvec
 from cutlass import apply_swizzle as _apply_smem_swizzle
@@ -74,33 +78,6 @@ def _auto_swizzle_w(m, n, k, nt_n):
     return cutlass.Int32(w)
 
 
-def _l2_swizzle_tile(raw_m, raw_n, nt_m, nt_n, swizzle_w):
-    """N-direction super-block rasterization of the (m, n) cgrp-tile coord, for
-    L2 reuse. ``swizzle_w == 1`` falls out of the math as the identity mapping.
-    """
-    t = raw_n * nt_m + raw_m
-    blk = nt_m * swizzle_w
-    sb = t // blk
-    off = t - sb * blk
-    base_n = sb * swizzle_w
-    cur_S = cutlass.min(cutlass.Int32(swizzle_w), nt_n - base_n)
-    log_m = off // cur_S
-    log_n = base_n + off - log_m * cur_S
-    return log_m, log_n
-
-
-def _epi_subtile_spans(cols):
-    spans = []
-    off = 0
-    while off < cols:
-        w = 32
-        while w > cols - off:
-            w //= 2
-        spans.append((off, w))
-        off += w
-    return spans
-
-
 @cute.kernel
 def _kernel(
     m: cutlass.Int64,
@@ -566,7 +543,12 @@ def _kernel(
 
     if warp_idx == mma_warp_id:
         nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
-        nvvm.tcgen05_alloc(tmem_ptr_i32, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_1)
+        nvvm.tcgen05_alloc(
+            tmem_ptr_i32,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_1,
+        )
         nvvm.bar_warp_sync(0xFFFFFFFF)
         nvvm.barrier_cta_arrive(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
         tmem_raw_addr = tmem_ptr_i32.load()
@@ -699,7 +681,12 @@ def _kernel(
         nvvm.bar_warp_sync(0xFFFFFFFF)
         nvvm.tcgen05_relinquish_alloc_permit(group=nvvm.CTAGroup.CTA_1)
         alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-        nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_1)
+        nvvm.tcgen05_dealloc(
+            alloc_ptr,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_1,
+        )
 
     if warp_idx < num_epilogue_warps:
         nvvm.setmaxregister(epi_reg_count, nvvm.SetMaxRegisterAction.INCREASE)
diff --git a/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_1ctamma_static.py b/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_1ctamma_static.py
index 21616928f..8f018f03d 100644
--- a/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_1ctamma_static.py
+++ b/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_1ctamma_static.py
@@ -23,6 +23,10 @@
 from typing import Callable
 
 import cutlass.experimental.primitives as nvvm
+from cudnn.gemm.frost.kernel_templates._tile_helpers import (
+    epi_subtile_spans as _epi_subtile_spans,
+    l2_swizzle_tile as _l2_swizzle_tile,
+)
 import cutlass.experimental.cuda.tensor_map as _tma
 import cutlass._mlir_helpers.vector as _cvec
 from cutlass import apply_swizzle as _apply_smem_swizzle
@@ -71,33 +75,6 @@ def _auto_swizzle_w(m, n, k, nt_n):
     return cutlass.Int32(w)
 
 
-def _l2_swizzle_tile(raw_m, raw_n, nt_m, nt_n, swizzle_w):
-    """N-direction super-block rasterization of the (m, n) cgrp-tile coord, for
-    L2 reuse. ``swizzle_w == 1`` falls out of the math as the identity mapping.
-    """
-    t = raw_n * nt_m + raw_m
-    blk = nt_m * swizzle_w
-    sb = t // blk
-    off = t - sb * blk
-    base_n = sb * swizzle_w
-    cur_S = cutlass.min(cutlass.Int32(swizzle_w), nt_n - base_n)
-    log_m = off // cur_S
-    log_n = base_n + off - log_m * cur_S
-    return log_m, log_n
-
-
-def _epi_subtile_spans(cols):
-    spans = []
-    off = 0
-    while off < cols:
-        w = 32
-        while w > cols - off:
-            w //= 2
-        spans.append((off, w))
-        off += w
-    return spans
-
-
 @cute.kernel
 def _kernel(
     m: cutlass.Int64,
@@ -463,7 +440,12 @@ def _kernel(
 
     if warp_idx == mma_warp_id:
         nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
-        nvvm.tcgen05_alloc(tmem_ptr_i32, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_1)
+        nvvm.tcgen05_alloc(
+            tmem_ptr_i32,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_1,
+        )
         nvvm.bar_warp_sync(0xFFFFFFFF)
         nvvm.barrier_cta_arrive(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
         tmem_raw_addr = tmem_ptr_i32.load()
@@ -577,7 +559,12 @@ def _kernel(
         nvvm.bar_warp_sync(0xFFFFFFFF)
         nvvm.tcgen05_relinquish_alloc_permit(group=nvvm.CTAGroup.CTA_1)
         alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-        nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_1)
+        nvvm.tcgen05_dealloc(
+            alloc_ptr,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_1,
+        )
 
     if warp_idx < num_epilogue_warps:
         nvvm.setmaxregister(epi_reg_count, nvvm.SetMaxRegisterAction.INCREASE)
diff --git a/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_2ctamma.py b/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_2ctamma.py
index 8f539da80..d69dd8256 100644
--- a/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_2ctamma.py
+++ b/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_2ctamma.py
@@ -22,6 +22,10 @@
 from typing import Callable
 
 import cutlass.experimental.primitives as nvvm
+from cudnn.gemm.frost.kernel_templates._tile_helpers import (
+    epi_subtile_spans as _epi_subtile_spans,
+    l2_swizzle_tile as _l2_swizzle_tile,
+)
 import cutlass.experimental.cuda.tensor_map as _tma
 import cutlass._mlir_helpers.vector as _cvec
 from cutlass import apply_swizzle as _apply_smem_swizzle
@@ -72,33 +76,6 @@ def _auto_swizzle_w(m, n, k, nt_n):
     return cutlass.Int32(w)
 
 
-def _l2_swizzle_tile(raw_m, raw_n, nt_m, nt_n, swizzle_w):
-    """N-direction super-block rasterization of the (m, n) cgrp-tile coord, for
-    L2 reuse. ``swizzle_w == 1`` falls out of the math as the identity mapping.
-    """
-    t = raw_n * nt_m + raw_m
-    blk = nt_m * swizzle_w
-    sb = t // blk
-    off = t - sb * blk
-    base_n = sb * swizzle_w
-    cur_S = cutlass.min(cutlass.Int32(swizzle_w), nt_n - base_n)
-    log_m = off // cur_S
-    log_n = base_n + off - log_m * cur_S
-    return log_m, log_n
-
-
-def _epi_subtile_spans(cols):
-    spans = []
-    off = 0
-    while off < cols:
-        w = 32
-        while w > cols - off:
-            w //= 2
-        spans.append((off, w))
-        off += w
-    return spans
-
-
 @cute.kernel
 def _kernel(
     m: cutlass.Int64,
@@ -586,7 +563,12 @@ def _kernel(
     ab_empty_arrive_mask = cutlass.Int16(a_part | b_part)
     if warp_idx == mma_warp_id:
         nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
-        nvvm.tcgen05_alloc(tmem_ptr_i32, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+        nvvm.tcgen05_alloc(
+            tmem_ptr_i32,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_2,
+        )
         nvvm.bar_warp_sync(0xFFFFFFFF)
         nvvm.barrier_cta_arrive(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
         tmem_raw_addr = tmem_ptr_i32.load()
@@ -735,7 +717,12 @@ def _kernel(
                 pass
             nvvm.mbarrier_arrive(peer_mbar, scope=nvvm.MemScope.CLUSTER, relaxed=True)
             alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-            nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+            nvvm.tcgen05_dealloc(
+                alloc_ptr,
+                cutlass.Int32(num_tmem_alloc_cols),
+                is_exclusive=tmem_alloc_exclusive,
+                group=nvvm.CTAGroup.CTA_2,
+            )
         else:
             tile_iter = cutlass.Int32(0)
             is_valid = cutlass.Int32(1)
@@ -769,7 +756,12 @@ def _kernel(
             while not nvvm.mbarrier_try_wait_parity(tmem_dealloc_mbar_ptr, 0, time_limit=10_000_000):
                 pass
             alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-            nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+            nvvm.tcgen05_dealloc(
+                alloc_ptr,
+                cutlass.Int32(num_tmem_alloc_cols),
+                is_exclusive=tmem_alloc_exclusive,
+                group=nvvm.CTAGroup.CTA_2,
+            )
 
     if warp_idx < num_epilogue_warps:
         nvvm.setmaxregister(epi_reg_count, nvvm.SetMaxRegisterAction.INCREASE)
diff --git a/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_2ctamma_static.py b/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_2ctamma_static.py
index 4baab1b2b..8c2c2e7c9 100644
--- a/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_2ctamma_static.py
+++ b/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_2ctamma_static.py
@@ -22,6 +22,10 @@
 from typing import Callable
 
 import cutlass.experimental.primitives as nvvm
+from cudnn.gemm.frost.kernel_templates._tile_helpers import (
+    epi_subtile_spans as _epi_subtile_spans,
+    l2_swizzle_tile as _l2_swizzle_tile,
+)
 import cutlass.experimental.cuda.tensor_map as _tma
 import cutlass._mlir_helpers.vector as _cvec
 from cutlass import apply_swizzle as _apply_smem_swizzle
@@ -69,33 +73,6 @@ def _auto_swizzle_w(m, n, k, nt_n):
     return cutlass.Int32(w)
 
 
-def _l2_swizzle_tile(raw_m, raw_n, nt_m, nt_n, swizzle_w):
-    """N-direction super-block rasterization of the (m, n) cgrp-tile coord, for
-    L2 reuse. ``swizzle_w == 1`` falls out of the math as the identity mapping.
-    """
-    t = raw_n * nt_m + raw_m
-    blk = nt_m * swizzle_w
-    sb = t // blk
-    off = t - sb * blk
-    base_n = sb * swizzle_w
-    cur_S = cutlass.min(cutlass.Int32(swizzle_w), nt_n - base_n)
-    log_m = off // cur_S
-    log_n = base_n + off - log_m * cur_S
-    return log_m, log_n
-
-
-def _epi_subtile_spans(cols):
-    spans = []
-    off = 0
-    while off < cols:
-        w = 32
-        while w > cols - off:
-            w //= 2
-        spans.append((off, w))
-        off += w
-    return spans
-
-
 @cute.kernel
 def _kernel(
     m: cutlass.Int64,
@@ -484,7 +461,12 @@ def _kernel(
     ab_empty_arrive_mask = cutlass.Int16(a_part | b_part)
     if warp_idx == mma_warp_id:
         nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
-        nvvm.tcgen05_alloc(tmem_ptr_i32, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+        nvvm.tcgen05_alloc(
+            tmem_ptr_i32,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_2,
+        )
         nvvm.bar_warp_sync(0xFFFFFFFF)
         nvvm.barrier_cta_arrive(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
         tmem_raw_addr = tmem_ptr_i32.load()
@@ -617,7 +599,12 @@ def _kernel(
                 pass
             nvvm.mbarrier_arrive(peer_mbar, scope=nvvm.MemScope.CLUSTER, relaxed=True)
             alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-            nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+            nvvm.tcgen05_dealloc(
+                alloc_ptr,
+                cutlass.Int32(num_tmem_alloc_cols),
+                is_exclusive=tmem_alloc_exclusive,
+                group=nvvm.CTAGroup.CTA_2,
+            )
         else:
             if cutlass.const_expr(USE_PDL):
                 if elect_one:
@@ -629,7 +616,12 @@ def _kernel(
             while not nvvm.mbarrier_try_wait_parity(tmem_dealloc_mbar_ptr, 0, time_limit=10_000_000):
                 pass
             alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-            nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+            nvvm.tcgen05_dealloc(
+                alloc_ptr,
+                cutlass.Int32(num_tmem_alloc_cols),
+                is_exclusive=tmem_alloc_exclusive,
+                group=nvvm.CTAGroup.CTA_2,
+            )
 
     if warp_idx < num_epilogue_warps:
         nvvm.setmaxregister(epi_reg_count, nvvm.SetMaxRegisterAction.INCREASE)
diff --git a/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_mainloop_1ctamma.py b/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_mainloop_1ctamma.py
index ec2400e79..c57d96fcb 100644
--- a/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_mainloop_1ctamma.py
+++ b/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_mainloop_1ctamma.py
@@ -23,6 +23,10 @@
 from typing import Callable
 
 import cutlass.experimental.primitives as nvvm
+from cudnn.gemm.frost.kernel_templates._tile_helpers import (
+    epi_subtile_spans as _epi_subtile_spans,
+    l2_swizzle_tile as _l2_swizzle_tile,
+)
 import cutlass.experimental.cuda.tensor_map as _tma
 import cutlass._mlir_helpers.vector as _cvec
 from cutlass import apply_swizzle as _apply_smem_swizzle
@@ -72,33 +76,6 @@ def _auto_swizzle_w(m, n, k, nt_n):
     return cutlass.Int32(w)
 
 
-def _l2_swizzle_tile(raw_m, raw_n, nt_m, nt_n, swizzle_w):
-    """N-direction super-block rasterization of the (m, n) cgrp-tile coord, for
-    L2 reuse. ``swizzle_w == 1`` falls out of the math as the identity mapping.
-    """
-    t = raw_n * nt_m + raw_m
-    blk = nt_m * swizzle_w
-    sb = t // blk
-    off = t - sb * blk
-    base_n = sb * swizzle_w
-    cur_S = cutlass.min(cutlass.Int32(swizzle_w), nt_n - base_n)
-    log_m = off // cur_S
-    log_n = base_n + off - log_m * cur_S
-    return log_m, log_n
-
-
-def _epi_subtile_spans(cols):
-    spans = []
-    off = 0
-    while off < cols:
-        w = 32
-        while w > cols - off:
-            w //= 2
-        spans.append((off, w))
-        off += w
-    return spans
-
-
 @cute.kernel
 def _kernel(
     m: cutlass.Int64,
@@ -579,7 +556,12 @@ def _kernel(
 
     if warp_idx == mma_warp_id:
         nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
-        nvvm.tcgen05_alloc(tmem_ptr_i32, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_1)
+        nvvm.tcgen05_alloc(
+            tmem_ptr_i32,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_1,
+        )
         nvvm.bar_warp_sync(0xFFFFFFFF)
         nvvm.barrier_cta_arrive(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
         tmem_raw_addr = tmem_ptr_i32.load()
@@ -725,7 +707,12 @@ def _kernel(
         nvvm.bar_warp_sync(0xFFFFFFFF)
         nvvm.tcgen05_relinquish_alloc_permit(group=nvvm.CTAGroup.CTA_1)
         alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-        nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_1)
+        nvvm.tcgen05_dealloc(
+            alloc_ptr,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_1,
+        )
 
     if warp_idx >= mainloop_warp_id_start:
         nvvm.setmaxregister(epi_reg_count, nvvm.SetMaxRegisterAction.INCREASE)
diff --git a/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_mainloop_2ctamma.py b/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_mainloop_2ctamma.py
index 32433eed0..5270ee597 100644
--- a/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_mainloop_2ctamma.py
+++ b/python/cudnn/gemm/frost/kernel_templates/sm100_matmul_mainloop_2ctamma.py
@@ -24,6 +24,10 @@
 from typing import Callable
 
 import cutlass.experimental.primitives as nvvm
+from cudnn.gemm.frost.kernel_templates._tile_helpers import (
+    epi_subtile_spans as _epi_subtile_spans,
+    l2_swizzle_tile as _l2_swizzle_tile,
+)
 import cutlass.experimental.cuda.tensor_map as _tma
 import cutlass._mlir_helpers.vector as _cvec
 from cutlass import apply_swizzle as _apply_smem_swizzle
@@ -74,33 +78,6 @@ def _auto_swizzle_w(m, n, k, nt_n):
     return cutlass.Int32(w)
 
 
-def _l2_swizzle_tile(raw_m, raw_n, nt_m, nt_n, swizzle_w):
-    """N-direction super-block rasterization of the (m, n) cgrp-tile coord, for
-    L2 reuse. ``swizzle_w == 1`` falls out of the math as the identity mapping.
-    """
-    t = raw_n * nt_m + raw_m
-    blk = nt_m * swizzle_w
-    sb = t // blk
-    off = t - sb * blk
-    base_n = sb * swizzle_w
-    cur_S = cutlass.min(cutlass.Int32(swizzle_w), nt_n - base_n)
-    log_m = off // cur_S
-    log_n = base_n + off - log_m * cur_S
-    return log_m, log_n
-
-
-def _epi_subtile_spans(cols):
-    spans = []
-    off = 0
-    while off < cols:
-        w = 32
-        while w > cols - off:
-            w //= 2
-        spans.append((off, w))
-        off += w
-    return spans
-
-
 @cute.kernel
 def _kernel(
     m: cutlass.Int64,
@@ -675,7 +652,12 @@ def _kernel(
     ab_empty_arrive_mask = cutlass.Int16(a_part | b_part)
     if warp_idx == mma_warp_id:
         nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
-        nvvm.tcgen05_alloc(tmem_ptr_i32, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+        nvvm.tcgen05_alloc(
+            tmem_ptr_i32,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_2,
+        )
         nvvm.bar_warp_sync(0xFFFFFFFF)
         nvvm.barrier_cta_arrive(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
         tmem_raw_addr = tmem_ptr_i32.load()
@@ -833,7 +815,12 @@ def _kernel(
                 pass
             nvvm.mbarrier_arrive(peer_mbar, scope=nvvm.MemScope.CLUSTER, relaxed=True)
             alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-            nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+            nvvm.tcgen05_dealloc(
+                alloc_ptr,
+                cutlass.Int32(num_tmem_alloc_cols),
+                is_exclusive=tmem_alloc_exclusive,
+                group=nvvm.CTAGroup.CTA_2,
+            )
         else:
             tile_iter = cutlass.Int32(0)
             is_valid = cutlass.Int32(1)
@@ -867,7 +854,12 @@ def _kernel(
             while not nvvm.mbarrier_try_wait_parity(tmem_dealloc_mbar_ptr, 0, time_limit=10_000_000):
                 pass
             alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-            nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+            nvvm.tcgen05_dealloc(
+                alloc_ptr,
+                cutlass.Int32(num_tmem_alloc_cols),
+                is_exclusive=tmem_alloc_exclusive,
+                group=nvvm.CTAGroup.CTA_2,
+            )
 
     if warp_idx >= mainloop_warp_id_start:
         nvvm.setmaxregister(epi_reg_count, nvvm.SetMaxRegisterAction.INCREASE)
diff --git a/python/cudnn/gemm/frost/kernel_templates/sm100_moe_grouped_block_scale_matmul_fwd_1ctamma.py b/python/cudnn/gemm/frost/kernel_templates/sm100_moe_grouped_block_scale_matmul_fwd_1ctamma.py
index a0cfedb64..404b91ffa 100644
--- a/python/cudnn/gemm/frost/kernel_templates/sm100_moe_grouped_block_scale_matmul_fwd_1ctamma.py
+++ b/python/cudnn/gemm/frost/kernel_templates/sm100_moe_grouped_block_scale_matmul_fwd_1ctamma.py
@@ -19,6 +19,15 @@
 from typing import Callable
 
 import cutlass.experimental.primitives as nvvm
+from cudnn.gemm.frost.kernel_templates._tile_helpers import (
+    copy_tensormap_to_workspace as _copy_tensormap_to_workspace,
+    fence_tensormap_acquire as _fence_tensormap_acquire,
+    fence_tensormap_release as _fence_tensormap_release,
+    moe_swizzle_tile as _moe_swizzle_tile,
+    replace_tensormap_global_address as _replace_tensormap_global_address,
+    replace_tensormap_global_dim_1 as _replace_tensormap_global_dim_1,
+    TENSOR_MAP_QWORDS,
+)
 import cutlass.experimental.cuda.tensor_map as _tma
 import cutlass
 import cutlass.cute as cute
@@ -29,56 +38,6 @@
 # A TMA tensormap is 128 bytes = 16 int64 qwords. The per-group A descriptor
 # replacement keeps a per-CTA SMEM copy, patches base/M-dim there, then publishes
 # it to the per-CTA GMEM workspace the TMA reads.
-_TENSOR_MAP_QWORDS = 16
-
-
-@cute.jit
-def _copy_tensormap_to_workspace(src_desc_ptr, dst_i64_ptr) -> None:
-    """Copy the 128-byte A tensormap into ``dst_i64_ptr`` (seeds the SMEM copy)."""
-    src_words = cute.make_ptr(cutlass.Int64, src_desc_ptr.toint(), mem_space=cute.AddressSpace.generic)
-    for i in cutlass.range_constexpr(_TENSOR_MAP_QWORDS):
-        dst_i64_ptr.subview(i).store((src_words + i).load())
-
-
-@cute.jit
-def _replace_tensormap_global_address(desc_ptr, new_address) -> None:
-    nvvm.tensormap_replace(
-        nvvm.TensormapField.GLOBAL_ADDRESS,
-        desc_ptr,
-        new_value=cutlass.Int64(new_address),
-    )
-
-
-@cute.jit
-def _replace_tensormap_global_dim_1(desc_ptr, new_dim) -> None:
-    nvvm.tensormap_replace(
-        nvvm.TensormapField.GLOBAL_DIM,
-        desc_ptr,
-        new_value=cutlass.Int32(new_dim),
-        ord=1,
-    )
-
-
-@cute.jit
-def _fence_tensormap_release() -> None:
-    nvvm.fence_proxy_release(
-        nvvm.MemScope.GPU,
-        from_proxy=nvvm.Proxy.GENERIC,
-        to_proxy=nvvm.Proxy.TENSORMAP,
-    )
-
-
-@cute.jit
-def _fence_tensormap_acquire(desc_ptr) -> None:
-    nvvm.fence_proxy_acquire(
-        nvvm.MemScope.GPU,
-        desc_ptr,
-        _TENSOR_MAP_QWORDS * 8,
-        from_proxy=nvvm.Proxy.GENERIC,
-        to_proxy=nvvm.Proxy.TENSORMAP,
-    )
-
-
 # @@INJECT_TILE_CONSTANTS@@
 
 
@@ -112,18 +71,18 @@ def _moe_auto_swizzle_w(group_rows, n, k, nt_n):
     return cutlass.Int32(w)
 
 
-def _moe_swizzle_tile(t, nt_m, nt_n, swizzle_w):
-    """Group-local linear tile index -> (m, n) under an N-super-block walk.
-    ``swizzle_w == nt_n`` reproduces the plain n-fast split; ``1`` gives m-fast.
-    """
-    blk = cutlass.max(nt_m * swizzle_w, cutlass.Int32(1))
-    sb = t // blk
-    off = t - sb * blk
-    base_n = sb * swizzle_w
-    cur_S = cutlass.min(cutlass.Int32(swizzle_w), nt_n - base_n)
-    tile_m = off // cur_S
-    tile_n = base_n + off - tile_m * cur_S
-    return tile_m, tile_n
+def _b_collector_op(mi):
+    """B is identical across the M sub-blocks (only A's address advances), so the
+    first MMA fills the B collector and the rest read it back instead of
+    re-fetching the same operand from SMEM. `.collector::b::*` is silicon-gated
+    (sm_107a only), hence `b_collector_ok`."""
+    if cutlass.const_expr(not b_collector_ok or num_mma_m == 1):
+        return None
+    if cutlass.const_expr(mi == 0):
+        return nvvm.Tcgen05MMACollectorOp.FILL
+    if cutlass.const_expr(mi == num_mma_m - 1):
+        return nvvm.Tcgen05MMACollectorOp.LASTUSE
+    return nvvm.Tcgen05MMACollectorOp.USE
 
 
 @cute.kernel
@@ -220,7 +179,7 @@ def _kernel(
     tma_a_desc_smem_list = [
         cutlass.Array(
             cutlass.Int64,
-            _TENSOR_MAP_QWORDS,
+            TENSOR_MAP_QWORDS,
             space=cutlass.AddressSpace.smem,
             alignment=128,
         )
@@ -541,7 +500,7 @@ def _kernel(
 
         lane = tidx % 32
         block_linear = bidx + bidy * gridx
-        cta_desc_base_list = [a_tma_workspace.iterator.raw_ptr() + (block_linear * num_a_operands + _ai) * _TENSOR_MAP_QWORDS for _ai in range(num_a_operands)]
+        cta_desc_base_list = [a_tma_workspace.iterator.raw_ptr() + (block_linear * num_a_operands + _ai) * TENSOR_MAP_QWORDS for _ai in range(num_a_operands)]
         a_desc_tma_ptr_list = [
             cute.make_ptr(
                 cutlass.Int64,
@@ -594,7 +553,7 @@ def _kernel(
                             _replace_tensormap_global_address(tma_a_desc_smem_list[_ai], row_base)
                             _replace_tensormap_global_dim_1(tma_a_desc_smem_list[_ai], group_end - group_begin)
                         nvvm.bar_warp_sync(0xFFFFFFFF)
-                        if lane < _TENSOR_MAP_QWORDS:
+                        if lane < TENSOR_MAP_QWORDS:
                             (cta_desc_base_list[_ai] + lane).store((tma_a_desc_smem_list[_ai].subview(lane)).load())
                         nvvm.bar_warp_sync(0xFFFFFFFF)
                         _fence_tensormap_release()
@@ -698,7 +657,12 @@ def _kernel(
 
     if warp_idx == mma_warp_id:
         nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
-        nvvm.tcgen05_alloc(tmem_ptr_i32, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_1)
+        nvvm.tcgen05_alloc(
+            tmem_ptr_i32,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_1,
+        )
         nvvm.bar_warp_sync(0xFFFFFFFF)
         nvvm.barrier_cta_arrive(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
         tmem_raw_addr = tmem_ptr_i32.load()
@@ -880,6 +844,7 @@ def _kernel(
                                             scale_a=sfa_dst_ptrs[_ai][mi],
                                             scale_b=sfb_scale_ptrs[_bj],
                                             scale_vec_size=scale_vec_size,
+                                            b_collector_op=_b_collector_op(mi),
                                         )
                             # Every accumulator sees scale_d=False on exactly the first
                             # k_block of the tile, so the flip stays outside mi.
@@ -926,7 +891,12 @@ def _kernel(
 
         nvvm.bar_warp_sync(0xFFFFFFFF)
         alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-        nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_1)
+        nvvm.tcgen05_dealloc(
+            alloc_ptr,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_1,
+        )
 
     if warp_idx < num_epilogue_warps:
         nvvm.setmaxregister(epi_reg_count, nvvm.SetMaxRegisterAction.INCREASE)
diff --git a/python/cudnn/gemm/frost/kernel_templates/sm100_moe_grouped_block_scale_matmul_fwd_2ctamma.py b/python/cudnn/gemm/frost/kernel_templates/sm100_moe_grouped_block_scale_matmul_fwd_2ctamma.py
index dbbcd3ed1..e7af83c91 100644
--- a/python/cudnn/gemm/frost/kernel_templates/sm100_moe_grouped_block_scale_matmul_fwd_2ctamma.py
+++ b/python/cudnn/gemm/frost/kernel_templates/sm100_moe_grouped_block_scale_matmul_fwd_2ctamma.py
@@ -24,6 +24,15 @@
 from typing import Callable
 
 import cutlass.experimental.primitives as nvvm
+from cudnn.gemm.frost.kernel_templates._tile_helpers import (
+    copy_tensormap_to_workspace as _copy_tensormap_to_workspace,
+    fence_tensormap_acquire as _fence_tensormap_acquire,
+    fence_tensormap_release as _fence_tensormap_release,
+    moe_swizzle_tile as _moe_swizzle_tile,
+    replace_tensormap_global_address as _replace_tensormap_global_address,
+    replace_tensormap_global_dim_1 as _replace_tensormap_global_dim_1,
+    TENSOR_MAP_QWORDS,
+)
 import cutlass.experimental.cuda.tensor_map as _tma
 import cutlass
 import cutlass.cute as cute
@@ -31,56 +40,6 @@
 from cutlass.cute.runtime import make_fake_stream
 from cuda.bindings import driver as _cuda
 
-_TENSOR_MAP_QWORDS = 16
-
-
-@cute.jit
-def _copy_tensormap_to_workspace(src_desc_ptr, dst_i64_ptr) -> None:
-    """Copy the 128-byte A tensormap into ``dst_i64_ptr`` (seeds the SMEM copy)."""
-    src_words = cute.make_ptr(cutlass.Int64, src_desc_ptr.toint(), mem_space=cute.AddressSpace.generic)
-    for i in cutlass.range(_TENSOR_MAP_QWORDS, unroll_full=True):
-        dst_i64_ptr.subview(i).store((src_words + i).load())
-
-
-@cute.jit
-def _replace_tensormap_global_address(desc_ptr, new_address) -> None:
-    nvvm.tensormap_replace(
-        nvvm.TensormapField.GLOBAL_ADDRESS,
-        desc_ptr,
-        new_value=cutlass.Int64(new_address),
-    )
-
-
-@cute.jit
-def _replace_tensormap_global_dim_1(desc_ptr, new_dim) -> None:
-    nvvm.tensormap_replace(
-        nvvm.TensormapField.GLOBAL_DIM,
-        desc_ptr,
-        new_value=cutlass.Int32(new_dim),
-        ord=1,
-    )
-
-
-@cute.jit
-def _fence_tensormap_release() -> None:
-    nvvm.fence_proxy_release(
-        nvvm.MemScope.GPU,
-        from_proxy=nvvm.Proxy.GENERIC,
-        to_proxy=nvvm.Proxy.TENSORMAP,
-    )
-
-
-@cute.jit
-def _fence_tensormap_acquire(desc_ptr) -> None:
-    nvvm.fence_proxy_acquire(
-        nvvm.MemScope.GPU,
-        desc_ptr,
-        _TENSOR_MAP_QWORDS * 8,
-        from_proxy=nvvm.Proxy.GENERIC,
-        to_proxy=nvvm.Proxy.TENSORMAP,
-    )
-
-
 # @@INJECT_TILE_CONSTANTS@@
 
 
@@ -113,18 +72,18 @@ def _moe_auto_swizzle_w(group_rows, n, k, nt_n):
     return cutlass.Int32(w)
 
 
-def _moe_swizzle_tile(t, nt_m, nt_n, swizzle_w):
-    """Group-local linear tile index -> (m, n) under an N-super-block walk.
-    ``swizzle_w == nt_n`` reproduces the plain n-fast split; ``1`` gives m-fast.
-    """
-    blk = cutlass.max(nt_m * swizzle_w, cutlass.Int32(1))
-    sb = t // blk
-    off = t - sb * blk
-    base_n = sb * swizzle_w
-    cur_S = cutlass.min(cutlass.Int32(swizzle_w), nt_n - base_n)
-    tile_m = off // cur_S
-    tile_n = base_n + off - tile_m * cur_S
-    return tile_m, tile_n
+def _b_collector_op(mi):
+    """B is identical across the M sub-blocks (only A's address advances), so the
+    first MMA fills the B collector and the rest read it back instead of
+    re-fetching the same operand from SMEM. `.collector::b::*` is silicon-gated
+    (sm_107a only), hence `b_collector_ok`."""
+    if cutlass.const_expr(not b_collector_ok or num_mma_m == 1):
+        return None
+    if cutlass.const_expr(mi == 0):
+        return nvvm.Tcgen05MMACollectorOp.FILL
+    if cutlass.const_expr(mi == num_mma_m - 1):
+        return nvvm.Tcgen05MMACollectorOp.LASTUSE
+    return nvvm.Tcgen05MMACollectorOp.USE
 
 
 @cute.kernel
@@ -223,7 +182,7 @@ def _kernel(
     tma_a_desc_smem_list = [
         cutlass.Array(
             cutlass.Int64,
-            _TENSOR_MAP_QWORDS,
+            TENSOR_MAP_QWORDS,
             space=cutlass.AddressSpace.smem,
             alignment=128,
         )
@@ -555,7 +514,7 @@ def _kernel(
 
         lane = tidx % 32
         block_linear = bidx + bidy * gridx
-        cta_desc_base_list = [a_tma_workspace.iterator.raw_ptr() + (block_linear * num_a_operands + _ai) * _TENSOR_MAP_QWORDS for _ai in range(num_a_operands)]
+        cta_desc_base_list = [a_tma_workspace.iterator.raw_ptr() + (block_linear * num_a_operands + _ai) * TENSOR_MAP_QWORDS for _ai in range(num_a_operands)]
         a_desc_tma_ptr_list = [
             cute.make_ptr(
                 cutlass.Int64,
@@ -609,7 +568,7 @@ def _kernel(
                             _replace_tensormap_global_address(tma_a_desc_smem_list[_ai], row_base)
                             _replace_tensormap_global_dim_1(tma_a_desc_smem_list[_ai], group_end - group_begin)
                         nvvm.bar_warp_sync(0xFFFFFFFF)
-                        if lane < _TENSOR_MAP_QWORDS:
+                        if lane < TENSOR_MAP_QWORDS:
                             (cta_desc_base_list[_ai] + lane).store((tma_a_desc_smem_list[_ai].subview(lane)).load())
                         nvvm.bar_warp_sync(0xFFFFFFFF)
                         _fence_tensormap_release()
@@ -727,7 +686,12 @@ def _kernel(
     ab_empty_arrive_mask = cutlass.Int16(a_part | b_part)
     if warp_idx == mma_warp_id:
         nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
-        nvvm.tcgen05_alloc(tmem_ptr_i32, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+        nvvm.tcgen05_alloc(
+            tmem_ptr_i32,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_2,
+        )
         nvvm.bar_warp_sync(0xFFFFFFFF)
         nvvm.barrier_cta_arrive(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
         tmem_raw_addr = tmem_ptr_i32.load()
@@ -909,6 +873,7 @@ def _kernel(
                                                 scale_a=sfa_dst_ptrs[_ai][mi],
                                                 scale_b=sfb_scale_ptrs[_bj],
                                                 scale_vec_size=scale_vec_size,
+                                                b_collector_op=_b_collector_op(mi),
                                             )
                                 # Every accumulator sees scale_d=False on exactly the first
                                 # k_block of the tile, so the flip stays outside mi.
@@ -958,7 +923,12 @@ def _kernel(
             if cutlass.const_expr(not use_acc_overlap):
                 nvvm.mbarrier_arrive(peer_mbar, scope=nvvm.MemScope.CLUSTER, relaxed=True)
             alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-            nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+            nvvm.tcgen05_dealloc(
+                alloc_ptr,
+                cutlass.Int32(num_tmem_alloc_cols),
+                is_exclusive=tmem_alloc_exclusive,
+                group=nvvm.CTAGroup.CTA_2,
+            )
         else:
             is_valid = cutlass.Int32(1)
             sched_stage = cutlass.Int32(0)
@@ -989,7 +959,12 @@ def _kernel(
             while not nvvm.mbarrier_try_wait_parity(tmem_dealloc_mbar_ptr, 0, time_limit=10_000_000):
                 pass
             alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-            nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+            nvvm.tcgen05_dealloc(
+                alloc_ptr,
+                cutlass.Int32(num_tmem_alloc_cols),
+                is_exclusive=tmem_alloc_exclusive,
+                group=nvvm.CTAGroup.CTA_2,
+            )
 
     if warp_idx < num_epilogue_warps:
         nvvm.setmaxregister(epi_reg_count, nvvm.SetMaxRegisterAction.INCREASE)
diff --git a/python/cudnn/gemm/frost/kernel_templates/sm100_moe_grouped_matmul_fwd_1ctamma.py b/python/cudnn/gemm/frost/kernel_templates/sm100_moe_grouped_matmul_fwd_1ctamma.py
index 3d25bc10a..3ae4d6bc9 100644
--- a/python/cudnn/gemm/frost/kernel_templates/sm100_moe_grouped_matmul_fwd_1ctamma.py
+++ b/python/cudnn/gemm/frost/kernel_templates/sm100_moe_grouped_matmul_fwd_1ctamma.py
@@ -24,6 +24,17 @@
 from typing import Callable
 
 import cutlass.experimental.primitives as nvvm
+from cudnn.gemm.frost.kernel_templates._tile_helpers import (
+    copy_tensormap_to_workspace as _copy_tensormap_to_workspace,
+    epi_subtile_spans as _epi_subtile_spans,
+    fence_tensormap_acquire as _fence_tensormap_acquire,
+    fence_tensormap_release as _fence_tensormap_release,
+    moe_group_at as _moe_group_at,
+    moe_swizzle_tile as _moe_swizzle_tile,
+    replace_tensormap_global_address as _replace_tensormap_global_address,
+    replace_tensormap_global_dim_1 as _replace_tensormap_global_dim_1,
+    TENSOR_MAP_QWORDS,
+)
 import cutlass.experimental.cuda.tensor_map as _tma
 import cutlass
 import cutlass.cute as cute
@@ -31,56 +42,6 @@
 from cutlass.cute.runtime import make_fake_stream
 from cuda.bindings import driver as _cuda
 
-_TENSOR_MAP_QWORDS = 16
-
-
-@cute.jit
-def _copy_tensormap_to_workspace(src_desc_ptr, dst_i64_ptr) -> None:
-    """Copy the 128-byte A tensormap into ``dst_i64_ptr`` (seeds the SMEM copy)."""
-    src_words = cute.make_ptr(cutlass.Int64, src_desc_ptr.toint(), mem_space=cute.AddressSpace.generic)
-    for i in cutlass.range(_TENSOR_MAP_QWORDS, unroll_full=True):
-        dst_i64_ptr.subview(i).store((src_words + i).load())
-
-
-@cute.jit
-def _replace_tensormap_global_address(desc_ptr, new_address) -> None:
-    nvvm.tensormap_replace(
-        nvvm.TensormapField.GLOBAL_ADDRESS,
-        desc_ptr,
-        new_value=cutlass.Int64(new_address),
-    )
-
-
-@cute.jit
-def _replace_tensormap_global_dim_1(desc_ptr, new_dim) -> None:
-    nvvm.tensormap_replace(
-        nvvm.TensormapField.GLOBAL_DIM,
-        desc_ptr,
-        new_value=cutlass.Int32(new_dim),
-        ord=1,
-    )
-
-
-@cute.jit
-def _fence_tensormap_release() -> None:
-    nvvm.fence_proxy_release(
-        nvvm.MemScope.GPU,
-        from_proxy=nvvm.Proxy.GENERIC,
-        to_proxy=nvvm.Proxy.TENSORMAP,
-    )
-
-
-@cute.jit
-def _fence_tensormap_acquire(desc_ptr) -> None:
-    nvvm.fence_proxy_acquire(
-        nvvm.MemScope.GPU,
-        desc_ptr,
-        _TENSOR_MAP_QWORDS * 8,
-        from_proxy=nvvm.Proxy.GENERIC,
-        to_proxy=nvvm.Proxy.TENSORMAP,
-    )
-
-
 # @@INJECT_TILE_CONSTANTS@@
 
 
@@ -100,21 +61,6 @@ def _fence_tensormap_acquire(desc_ptr) -> None:
 TMEM_ALLOC_BARRIER_ID = 2
 
 
-@cute.jit
-def _moe_group_at(visit_idx, num_groups, num_experts):
-    """Visitation index -> routed group index.
-
-    ``num_groups == num_experts`` (or a non-multiple) walks groups in order. Batched MoE
-    (``num_groups == B * num_experts``) walks expert-major -- the B groups sharing expert
-    ``g % E`` become consecutive, so the expert weight is fetched once instead of B times.
-    """
-    per_expert = num_groups // cutlass.max(num_experts, cutlass.Int32(1))
-    group = visit_idx
-    if per_expert > 1 and per_expert * num_experts == num_groups:
-        group = (visit_idx % per_expert) * num_experts + (visit_idx // per_expert)
-    return group
-
-
 @cute.jit
 def _moe_auto_swizzle_w(group_rows, n, k, nt_n):
     """N-super-block width for one routed group, resolved per group.
@@ -135,32 +81,6 @@ def _moe_auto_swizzle_w(group_rows, n, k, nt_n):
     return cutlass.Int32(w)
 
 
-def _moe_swizzle_tile(t, nt_m, nt_n, swizzle_w):
-    """Group-local linear tile index -> (m, n) under an N-super-block walk.
-    ``swizzle_w == nt_n`` reproduces the plain n-fast split; ``1`` gives m-fast.
-    """
-    blk = cutlass.max(nt_m * swizzle_w, cutlass.Int32(1))
-    sb = t // blk
-    off = t - sb * blk
-    base_n = sb * swizzle_w
-    cur_S = cutlass.min(cutlass.Int32(swizzle_w), nt_n - base_n)
-    tile_m = off // cur_S
-    tile_n = base_n + off - tile_m * cur_S
-    return tile_m, tile_n
-
-
-def _epi_subtile_spans(cols):
-    spans = []
-    off = 0
-    while off < cols:
-        w = 32
-        while w > cols - off:
-            w //= 2
-        spans.append((off, w))
-        off += w
-    return spans
-
-
 @cute.kernel
 def _kernel(
     m: cutlass.Int64,
@@ -273,7 +193,7 @@ def _kernel(
     tma_a_desc_smem_list = [
         cutlass.Array(
             cutlass.Int64,
-            _TENSOR_MAP_QWORDS,
+            TENSOR_MAP_QWORDS,
             space=cutlass.AddressSpace.smem,
             alignment=128,
         )
@@ -463,7 +383,7 @@ def _kernel(
 
         lane = tidx % 32
         block_linear = bidx + bidy * gridx
-        cta_desc_base_list = [a_tma_workspace.iterator.raw_ptr() + (block_linear * num_a_operands + _ai) * _TENSOR_MAP_QWORDS for _ai in range(num_a_operands)]
+        cta_desc_base_list = [a_tma_workspace.iterator.raw_ptr() + (block_linear * num_a_operands + _ai) * TENSOR_MAP_QWORDS for _ai in range(num_a_operands)]
         a_desc_tma_ptr_list = [
             cute.make_ptr(
                 cutlass.Int64,
@@ -513,7 +433,7 @@ def _kernel(
                             _replace_tensormap_global_address(tma_a_desc_smem_list[_ai], row_base.toint())
                             _replace_tensormap_global_dim_1(tma_a_desc_smem_list[_ai], group_end - group_begin)
                         nvvm.bar_warp_sync(0xFFFFFFFF)
-                        if lane < _TENSOR_MAP_QWORDS:
+                        if lane < TENSOR_MAP_QWORDS:
                             (cta_desc_base_list[_ai] + lane).store((tma_a_desc_smem_list[_ai].subview(lane)).load())
                         nvvm.bar_warp_sync(0xFFFFFFFF)
                         _fence_tensormap_release()
@@ -596,7 +516,12 @@ def _kernel(
 
     if warp_idx == mma_warp_id:
         nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
-        nvvm.tcgen05_alloc(tmem_ptr_i32, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_1)
+        nvvm.tcgen05_alloc(
+            tmem_ptr_i32,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_1,
+        )
         nvvm.bar_warp_sync(0xFFFFFFFF)
         nvvm.barrier_cta_arrive(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
         tmem_raw_addr = tmem_ptr_i32.load()
@@ -737,7 +662,12 @@ def _kernel(
         nvvm.bar_warp_sync(0xFFFFFFFF)
         nvvm.tcgen05_relinquish_alloc_permit(group=nvvm.CTAGroup.CTA_1)
         alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-        nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_1)
+        nvvm.tcgen05_dealloc(
+            alloc_ptr,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_1,
+        )
 
     if warp_idx < num_epilogue_warps:
         nvvm.setmaxregister(epi_reg_count, nvvm.SetMaxRegisterAction.INCREASE)
diff --git a/python/cudnn/gemm/frost/kernel_templates/sm100_moe_grouped_matmul_fwd_2ctamma.py b/python/cudnn/gemm/frost/kernel_templates/sm100_moe_grouped_matmul_fwd_2ctamma.py
index 372aa6cb9..cc64ecc4a 100644
--- a/python/cudnn/gemm/frost/kernel_templates/sm100_moe_grouped_matmul_fwd_2ctamma.py
+++ b/python/cudnn/gemm/frost/kernel_templates/sm100_moe_grouped_matmul_fwd_2ctamma.py
@@ -25,6 +25,17 @@
 from typing import Callable
 
 import cutlass.experimental.primitives as nvvm
+from cudnn.gemm.frost.kernel_templates._tile_helpers import (
+    copy_tensormap_to_workspace as _copy_tensormap_to_workspace,
+    epi_subtile_spans as _epi_subtile_spans,
+    fence_tensormap_acquire as _fence_tensormap_acquire,
+    fence_tensormap_release as _fence_tensormap_release,
+    moe_group_at as _moe_group_at,
+    moe_swizzle_tile as _moe_swizzle_tile,
+    replace_tensormap_global_address as _replace_tensormap_global_address,
+    replace_tensormap_global_dim_1 as _replace_tensormap_global_dim_1,
+    TENSOR_MAP_QWORDS,
+)
 import cutlass.experimental.cuda.tensor_map as _tma
 import cutlass
 import cutlass.cute as cute
@@ -33,56 +44,6 @@
 from cuda.bindings import driver as _cuda
 
 # A TMA tensormap is 128 bytes = 16 int64 qwords.
-_TENSOR_MAP_QWORDS = 16
-
-
-@cute.jit
-def _copy_tensormap_to_workspace(src_desc_ptr, dst_i64_ptr) -> None:
-    """Copy the 128-byte A tensormap (``src_desc_ptr``) into ``dst_i64_ptr``."""
-    src_words = cute.make_ptr(cutlass.Int64, src_desc_ptr.toint(), mem_space=cute.AddressSpace.generic)
-    for i in cutlass.range(_TENSOR_MAP_QWORDS, unroll_full=True):
-        dst_i64_ptr.subview(i).store((src_words + i).load())
-
-
-@cute.jit
-def _replace_tensormap_global_address(desc_ptr, new_address) -> None:
-    nvvm.tensormap_replace(
-        nvvm.TensormapField.GLOBAL_ADDRESS,
-        desc_ptr,
-        new_value=cutlass.Int64(new_address),
-    )
-
-
-@cute.jit
-def _replace_tensormap_global_dim_1(desc_ptr, new_dim) -> None:
-    nvvm.tensormap_replace(
-        nvvm.TensormapField.GLOBAL_DIM,
-        desc_ptr,
-        new_value=cutlass.Int32(new_dim),
-        ord=1,
-    )
-
-
-@cute.jit
-def _fence_tensormap_release() -> None:
-    nvvm.fence_proxy_release(
-        nvvm.MemScope.GPU,
-        from_proxy=nvvm.Proxy.GENERIC,
-        to_proxy=nvvm.Proxy.TENSORMAP,
-    )
-
-
-@cute.jit
-def _fence_tensormap_acquire(desc_ptr) -> None:
-    nvvm.fence_proxy_acquire(
-        nvvm.MemScope.GPU,
-        desc_ptr,
-        _TENSOR_MAP_QWORDS * 8,
-        from_proxy=nvvm.Proxy.GENERIC,
-        to_proxy=nvvm.Proxy.TENSORMAP,
-    )
-
-
 # @@INJECT_TILE_CONSTANTS@@
 
 
@@ -105,21 +66,6 @@ def _fence_tensormap_acquire(desc_ptr) -> None:
 TMEM_ALLOC_BARRIER_ID = 2
 
 
-@cute.jit
-def _moe_group_at(visit_idx, num_groups, num_experts):
-    """Visitation index -> routed group index.
-
-    ``num_groups == num_experts`` (or a non-multiple) walks groups in order. Batched MoE
-    (``num_groups == B * num_experts``) walks expert-major -- the B groups sharing expert
-    ``g % E`` become consecutive, so the expert weight is fetched once instead of B times.
-    """
-    per_expert = num_groups // cutlass.max(num_experts, cutlass.Int32(1))
-    group = visit_idx
-    if per_expert > 1 and per_expert * num_experts == num_groups:
-        group = (visit_idx % per_expert) * num_experts + (visit_idx // per_expert)
-    return group
-
-
 @cute.jit
 def _moe_auto_swizzle_w(group_rows, n, k, nt_n):
     """N-super-block width for one routed group, resolved per group.
@@ -140,32 +86,6 @@ def _moe_auto_swizzle_w(group_rows, n, k, nt_n):
     return cutlass.Int32(w)
 
 
-def _moe_swizzle_tile(t, nt_m, nt_n, swizzle_w):
-    """Group-local linear tile index -> (m, n) under an N-super-block walk.
-    ``swizzle_w == nt_n`` reproduces the plain n-fast split; ``1`` gives m-fast.
-    """
-    blk = cutlass.max(nt_m * swizzle_w, cutlass.Int32(1))
-    sb = t // blk
-    off = t - sb * blk
-    base_n = sb * swizzle_w
-    cur_S = cutlass.min(cutlass.Int32(swizzle_w), nt_n - base_n)
-    tile_m = off // cur_S
-    tile_n = base_n + off - tile_m * cur_S
-    return tile_m, tile_n
-
-
-def _epi_subtile_spans(cols):
-    spans = []
-    off = 0
-    while off < cols:
-        w = 32
-        while w > cols - off:
-            w //= 2
-        spans.append((off, w))
-        off += w
-    return spans
-
-
 @cute.kernel
 def _kernel(
     m: cutlass.Int64,
@@ -281,7 +201,7 @@ def _kernel(
     tma_a_desc_smem_list = [
         cutlass.Array(
             cutlass.Int64,
-            _TENSOR_MAP_QWORDS,
+            TENSOR_MAP_QWORDS,
             space=cutlass.AddressSpace.smem,
             alignment=128,
         )
@@ -481,7 +401,7 @@ def _kernel(
 
         lane = tidx % 32
         block_linear = bidx + bidy * gridx
-        cta_desc_base_list = [a_tma_workspace.iterator.raw_ptr() + (block_linear * num_a_operands + _ai) * _TENSOR_MAP_QWORDS for _ai in range(num_a_operands)]
+        cta_desc_base_list = [a_tma_workspace.iterator.raw_ptr() + (block_linear * num_a_operands + _ai) * TENSOR_MAP_QWORDS for _ai in range(num_a_operands)]
         a_desc_tma_ptr_list = [
             cute.make_ptr(
                 cutlass.Int64,
@@ -531,7 +451,7 @@ def _kernel(
                             _replace_tensormap_global_address(tma_a_desc_smem_list[_ai], row_base.toint())
                             _replace_tensormap_global_dim_1(tma_a_desc_smem_list[_ai], group_end - group_begin)
                         nvvm.bar_warp_sync(0xFFFFFFFF)
-                        if lane < _TENSOR_MAP_QWORDS:
+                        if lane < TENSOR_MAP_QWORDS:
                             (cta_desc_base_list[_ai] + lane).store((tma_a_desc_smem_list[_ai].subview(lane)).load())
                         nvvm.bar_warp_sync(0xFFFFFFFF)
                         _fence_tensormap_release()
@@ -626,7 +546,12 @@ def _kernel(
     ab_empty_arrive_mask = cutlass.Int16(a_part | b_part)
     if warp_idx == mma_warp_id:
         nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
-        nvvm.tcgen05_alloc(tmem_ptr_i32, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+        nvvm.tcgen05_alloc(
+            tmem_ptr_i32,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_2,
+        )
         nvvm.bar_warp_sync(0xFFFFFFFF)
         nvvm.barrier_cta_arrive(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
         tmem_raw_addr = tmem_ptr_i32.load()
@@ -774,7 +699,12 @@ def _kernel(
                 pass
             nvvm.mbarrier_arrive(peer_mbar, scope=nvvm.MemScope.CLUSTER, relaxed=True)
             alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-            nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+            nvvm.tcgen05_dealloc(
+                alloc_ptr,
+                cutlass.Int32(num_tmem_alloc_cols),
+                is_exclusive=tmem_alloc_exclusive,
+                group=nvvm.CTAGroup.CTA_2,
+            )
         else:
             is_valid = cutlass.Int32(1)
             sched_stage = cutlass.Int32(0)
@@ -804,7 +734,12 @@ def _kernel(
             while not nvvm.mbarrier_try_wait_parity(tmem_dealloc_mbar_ptr, 0, time_limit=10_000_000):
                 pass
             alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-            nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+            nvvm.tcgen05_dealloc(
+                alloc_ptr,
+                cutlass.Int32(num_tmem_alloc_cols),
+                is_exclusive=tmem_alloc_exclusive,
+                group=nvvm.CTAGroup.CTA_2,
+            )
 
     if warp_idx < num_epilogue_warps:
         nvvm.setmaxregister(epi_reg_count, nvvm.SetMaxRegisterAction.INCREASE)
diff --git a/python/cudnn/gemm/frost/kernel_templates/sm103_block_scale_matmul_1ctamma.py b/python/cudnn/gemm/frost/kernel_templates/sm103_block_scale_matmul_1ctamma.py
index 5a20b19a2..65402c61f 100644
--- a/python/cudnn/gemm/frost/kernel_templates/sm103_block_scale_matmul_1ctamma.py
+++ b/python/cudnn/gemm/frost/kernel_templates/sm103_block_scale_matmul_1ctamma.py
@@ -5,7 +5,8 @@
 
 Computes ``C = (descale_a ⊙ A) @ (descale_b ⊙ B)`` where A/B are FP4 e2m1
 (packed 2-per-byte), dequantized by a per-block scale factor along K inside the
-MMA. FP4-only pipeline: nvfp4 (fp4/e4m3/block16) and mxfp4 (fp4/e8m0/block32).
+MMA. FP4-only pipeline: any of the e4m3 / e8m0 / e5m3 scale dtypes at either
+K-block (the two axes are orthogonal; e5m3 and e4m3-at-block-32 need SM 10.7+).
 
 sm103's fp4 UTCOMMA instruction K width is 48 BYTES (96 fp4 elements) vs
 sm100's 32 B (64 elements). 48 does not divide the 128-B swizzled SMEM line, so
@@ -43,6 +44,9 @@
 from typing import Callable
 
 import cutlass.experimental.primitives as nvvm
+from cudnn.gemm.frost.kernel_templates._tile_helpers import (
+    l2_swizzle_tile as _l2_swizzle_tile,
+)
 import cutlass.experimental.cuda.tensor_map as _tma
 import cutlass._mlir_helpers.vector as _cvec
 from cutlass import apply_swizzle as _apply_smem_swizzle
@@ -117,21 +121,6 @@ def _auto_swizzle_w(m, n, k, nt_n):
     return cutlass.Int32(w)
 
 
-def _l2_swizzle_tile(raw_m, raw_n, nt_m, nt_n, swizzle_w):
-    """N-direction super-block rasterization of the (m, n) cgrp-tile coord, for
-    L2 reuse. ``swizzle_w == 1`` falls out of the math as the identity mapping.
-    """
-    t = raw_n * nt_m + raw_m
-    blk = nt_m * swizzle_w
-    sb = t // blk
-    off = t - sb * blk
-    base_n = sb * swizzle_w
-    cur_S = cutlass.min(cutlass.Int32(swizzle_w), nt_n - base_n)
-    log_m = off // cur_S
-    log_n = base_n + off - log_m * cur_S
-    return log_m, log_n
-
-
 def _sm103_circular_mma_desc_base(current_desc):
     """Precompute invariant fields of an SM103 K=96 circular MMA SMEM desc:
     set the circular leading-dim mode bit (52) and clear the next-chunk
@@ -150,6 +139,20 @@ def _sm103_make_circular_mma_desc(current_desc_circular, phase_k16, next_addr_bi
     return nvvm.Tcgen05SmemDesc(desc_with_phase | next_addr_bits)
 
 
+def _b_collector_op(mi):
+    """B is identical across the M sub-blocks (only A's address advances), so the
+    first MMA fills the B collector and the rest read it back instead of
+    re-fetching the same operand from SMEM. `.collector::b::*` is silicon-gated
+    (sm_107a only), hence `b_collector_ok`."""
+    if cutlass.const_expr(not b_collector_ok or num_mma_m == 1):
+        return None
+    if cutlass.const_expr(mi == 0):
+        return nvvm.Tcgen05MMACollectorOp.FILL
+    if cutlass.const_expr(mi == num_mma_m - 1):
+        return nvvm.Tcgen05MMACollectorOp.LASTUSE
+    return nvvm.Tcgen05MMACollectorOp.USE
+
+
 @cute.kernel
 def _kernel(
     m: cutlass.Int64,
@@ -691,7 +694,12 @@ def _kernel(
 
     if warp_idx == mma_warp_id:
         nvvm.setmaxregister(mma_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
-        nvvm.tcgen05_alloc(tmem_ptr_i32, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_1)
+        nvvm.tcgen05_alloc(
+            tmem_ptr_i32,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_1,
+        )
         nvvm.bar_warp_sync(0xFFFFFFFF)
         nvvm.barrier_cta_arrive(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
         tmem_raw_addr = tmem_ptr_i32.load()
@@ -854,6 +862,7 @@ def _kernel(
                                         scale_a=nvvm.make_tmem_ptr(sfa_tmem_bases[_ai] + sfa_mma_col_off_by_j[_pj] + mi * registers_per_atom, cutlass.Float32),
                                         scale_b=nvvm.make_tmem_ptr(sfb_tmem_bases[_bj] + sfb_mma_col_off_by_j[_pj], cutlass.Float32),
                                         scale_vec_size=scale_vec_size,
+                                        b_collector_op=_b_collector_op(mi),
                                     )
                         scale_d = cutlass.Boolean(True)
                         for _rs in _RELEASE_SLOTS_AT.get(_pj, []):
@@ -964,6 +973,7 @@ def _kernel(
                                 scale_a=nvvm.make_tmem_ptr(sfa_tmem_bases[_ai] + sfa_mma_col_off_by_j[_pj] + mi * registers_per_atom, cutlass.Float32),
                                 scale_b=nvvm.make_tmem_ptr(sfb_tmem_bases[_bj] + sfb_mma_col_off_by_j[_pj], cutlass.Float32),
                                 scale_vec_size=scale_vec_size,
+                                b_collector_op=_b_collector_op(mi),
                             )
                 scale_d = cutlass.Boolean(True)
                 for _rs in _RELEASE_SLOTS_AT.get(_pj, []):
@@ -1022,7 +1032,12 @@ def _kernel(
 
         nvvm.bar_warp_sync(0xFFFFFFFF)
         alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-        nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_1)
+        nvvm.tcgen05_dealloc(
+            alloc_ptr,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_1,
+        )
 
     if warp_idx < num_epilogue_warps:
         nvvm.setmaxregister(epi_reg_count, nvvm.SetMaxRegisterAction.INCREASE)
diff --git a/python/cudnn/gemm/frost/kernel_templates/sm103_block_scale_matmul_2ctamma.py b/python/cudnn/gemm/frost/kernel_templates/sm103_block_scale_matmul_2ctamma.py
index 1a4b306d6..1974b5e11 100644
--- a/python/cudnn/gemm/frost/kernel_templates/sm103_block_scale_matmul_2ctamma.py
+++ b/python/cudnn/gemm/frost/kernel_templates/sm103_block_scale_matmul_2ctamma.py
@@ -5,7 +5,8 @@
 
 Computes ``C = (descale_a ⊙ A) @ (descale_b ⊙ B)`` where A/B are FP4 e2m1
 (packed 2-per-byte), dequantized by a per-block scale factor along K inside the
-MMA. FP4-only pipeline: nvfp4 (fp4/e4m3/block16) and mxfp4 (fp4/e8m0/block32).
+MMA. FP4-only pipeline: any of the e4m3 / e8m0 / e5m3 scale dtypes at either
+K-block (the two axes are orthogonal; e5m3 and e4m3-at-block-32 need SM 10.7+).
 
 sm103's fp4 UTCOMMA instruction K width is 48 BYTES (96 fp4 elements) vs
 sm100's 32 B (64 elements). 48 does not divide the 128-B swizzled SMEM line, so
@@ -43,6 +44,9 @@
 from typing import Callable
 
 import cutlass.experimental.primitives as nvvm
+from cudnn.gemm.frost.kernel_templates._tile_helpers import (
+    l2_swizzle_tile as _l2_swizzle_tile,
+)
 import cutlass.experimental.cuda.tensor_map as _tma
 import cutlass._mlir_helpers.vector as _cvec
 from cutlass import apply_swizzle as _apply_smem_swizzle
@@ -117,21 +121,6 @@ def _auto_swizzle_w(m, n, k, nt_n):
     return cutlass.Int32(w)
 
 
-def _l2_swizzle_tile(raw_m, raw_n, nt_m, nt_n, swizzle_w):
-    """N-direction super-block rasterization of the (m, n) cgrp-tile coord, for
-    L2 reuse. ``swizzle_w == 1`` falls out of the math as the identity mapping.
-    """
-    t = raw_n * nt_m + raw_m
-    blk = nt_m * swizzle_w
-    sb = t // blk
-    off = t - sb * blk
-    base_n = sb * swizzle_w
-    cur_S = cutlass.min(cutlass.Int32(swizzle_w), nt_n - base_n)
-    log_m = off // cur_S
-    log_n = base_n + off - log_m * cur_S
-    return log_m, log_n
-
-
 def _sm103_circular_mma_desc_base(current_desc):
     """Precompute invariant fields of an SM103 K=96 circular MMA SMEM desc:
     set the circular leading-dim mode bit (52) and clear the next-chunk
@@ -150,6 +139,20 @@ def _sm103_make_circular_mma_desc(current_desc_circular, phase_k16, next_addr_bi
     return nvvm.Tcgen05SmemDesc(desc_with_phase | next_addr_bits)
 
 
+def _b_collector_op(mi):
+    """B is identical across the M sub-blocks (only A's address advances), so the
+    first MMA fills the B collector and the rest read it back instead of
+    re-fetching the same operand from SMEM. `.collector::b::*` is silicon-gated
+    (sm_107a only), hence `b_collector_ok`."""
+    if cutlass.const_expr(not b_collector_ok or num_mma_m == 1):
+        return None
+    if cutlass.const_expr(mi == 0):
+        return nvvm.Tcgen05MMACollectorOp.FILL
+    if cutlass.const_expr(mi == num_mma_m - 1):
+        return nvvm.Tcgen05MMACollectorOp.LASTUSE
+    return nvvm.Tcgen05MMACollectorOp.USE
+
+
 @cute.kernel
 def _kernel(
     m: cutlass.Int64,
@@ -710,7 +713,12 @@ def _kernel(
 
     if warp_idx == mma_warp_id:
         nvvm.setmaxregister(mma_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
-        nvvm.tcgen05_alloc(tmem_ptr_i32, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+        nvvm.tcgen05_alloc(
+            tmem_ptr_i32,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_2,
+        )
         nvvm.bar_warp_sync(0xFFFFFFFF)
         nvvm.barrier_cta_arrive(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
         tmem_raw_addr = tmem_ptr_i32.load()
@@ -888,6 +896,7 @@ def _kernel(
                                             ),
                                             scale_b=nvvm.make_tmem_ptr(sfb_tmem_bases[_bj] + sfb_mma_col_off_by_j[_pj], cutlass.Float32),
                                             scale_vec_size=scale_vec_size,
+                                            b_collector_op=_b_collector_op(mi),
                                         )
                             scale_d = cutlass.Boolean(True)
                             for _rs in _RELEASE_SLOTS_AT.get(_pj, []):
@@ -998,6 +1007,7 @@ def _kernel(
                                     scale_a=nvvm.make_tmem_ptr(sfa_tmem_bases[_ai] + sfa_mma_col_off_by_j[_pj] + mi * registers_per_atom, cutlass.Float32),
                                     scale_b=nvvm.make_tmem_ptr(sfb_tmem_bases[_bj] + sfb_mma_col_off_by_j[_pj], cutlass.Float32),
                                     scale_vec_size=scale_vec_size,
+                                    b_collector_op=_b_collector_op(mi),
                                 )
                     scale_d = cutlass.Boolean(True)
                     for _rs in _RELEASE_SLOTS_AT.get(_pj, []):
@@ -1063,7 +1073,12 @@ def _kernel(
             if cutlass.const_expr(not use_acc_overlap):
                 nvvm.mbarrier_arrive(peer_mbar, scope=nvvm.MemScope.CLUSTER, relaxed=True)
             alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-            nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+            nvvm.tcgen05_dealloc(
+                alloc_ptr,
+                cutlass.Int32(num_tmem_alloc_cols),
+                is_exclusive=tmem_alloc_exclusive,
+                group=nvvm.CTAGroup.CTA_2,
+            )
         else:
             tile_iter = cutlass.Int32(0)
             is_valid = cutlass.Int32(1)
@@ -1098,7 +1113,12 @@ def _kernel(
             while not nvvm.mbarrier_try_wait_parity(tmem_dealloc_mbar_ptr, 0, time_limit=10_000_000):
                 pass
             alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
-            nvvm.tcgen05_dealloc(alloc_ptr, num_tmem_alloc_cols, group=nvvm.CTAGroup.CTA_2)
+            nvvm.tcgen05_dealloc(
+                alloc_ptr,
+                cutlass.Int32(num_tmem_alloc_cols),
+                is_exclusive=tmem_alloc_exclusive,
+                group=nvvm.CTAGroup.CTA_2,
+            )
 
     if warp_idx < num_epilogue_warps:
         nvvm.setmaxregister(epi_reg_count, nvvm.SetMaxRegisterAction.INCREASE)
diff --git a/python/cudnn/gemm/frost/kernel_templates/sm107_block_scale_matmul_1ctamma.py b/python/cudnn/gemm/frost/kernel_templates/sm107_block_scale_matmul_1ctamma.py
new file mode 100644
index 000000000..4569c960d
--- /dev/null
+++ b/python/cudnn/gemm/frost/kernel_templates/sm107_block_scale_matmul_1ctamma.py
@@ -0,0 +1,1442 @@
+# Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: MIT
+
+"""sm107 CTA_1 **block-scaled** GEMM kernel: persistent + CLC dynamic scheduler.
+
+Computes ``C = (descale_a ⊙ A) @ (descale_b ⊙ B)`` where A/B are narrow (FP4
+e2m1 packed 2-per-byte, or FP8 e4m3/e5m2), dequantized by a per-block scale
+factor along K inside the MMA. Supports nvfp4 (fp4/e4m3/block16), mxfp4
+(fp4/e8m0/block32), and mxfp8 (fp8/e8m0/block32). Single-CTA MMA; the compiler
+picks this when ``TileConfig.cta_group == 1``.
+
+The pipeline is the sm100 one; SM 10.7's block-scale MMA reads a **64-byte K**
+per instruction instead of 32, which shows up in exactly two places (both
+driven by injected constants, so the rest of the file stays in lockstep with
+``sm100_block_scale_matmul_1ctamma.py``):
+
+  * half as many MMAs per K-tile, each consuming ``sf_scales_per_inst`` scales
+    (8 at K-block 16, 4 at 32 — it follows the BLOCK SIZE, not the scale
+    dtype). When that exceeds the 4 scales one 128x4 utccp atom holds, a scale
+    *word* spans ``word_atoms`` atoms, and the two SF regions then lay them
+    out DIFFERENTLY: SFB atom-major across its N-blocks, SFA block-major.
+    At K-block 32 ``word_atoms == 1`` (identical to sm100).
+  * fp4 rides the OMMA instruction descriptor (``Tcgen05MxOmmaInstrDesc``,
+    K-mode 2 = 128 fp4 elements); mxfp8 stays on ``Tcgen05MxInstrDesc``
+    (K-mode 1 = 64 fp8 elements). Both take the real operand dtype.
+
+(The 576-column TMEM this GPU has is an ARCH property, not a pipeline one —
+`_TMEM_COLS_BY_ARCH` hands it to the sm100 templates on the same part too.)
+
+Warp layout (8 warps × 32 = 256 threads/CTA):
+  warps 0–3 : epilogue (warp 0 also allocates TMEM)  — setmaxnreg.inc 216
+  warp  4   : MMA driver (every CTA runs MMA — no pair structure)  — setmaxnreg.dec 40
+  warp  5   : TMA producer  — setmaxnreg.dec 40
+  warp  6   : CLC scheduler (leader CTA issues queries; every CTA waits + reads + arrives empty)  — setmaxnreg.dec 40
+  warp  7   : unused donor — setmaxnreg.dec 40, idle to dealloc barrier
+"""
+
+from __future__ import annotations
+
+from functools import lru_cache
+from typing import Callable
+
+import cutlass.experimental.primitives as nvvm
+from cudnn.gemm.frost.kernel_templates._tile_helpers import (
+    l2_swizzle_tile as _l2_swizzle_tile,
+)
+import cutlass.experimental.cuda.tensor_map as _tma
+import cutlass._mlir_helpers.vector as _cvec
+from cutlass import apply_swizzle as _apply_smem_swizzle
+import cutlass
+import cutlass.cute as cute
+from cutlass.cute.runtime import make_fake_compact_tensor
+from cutlass.cute.runtime import make_fake_stream
+from cuda.bindings import driver as _cuda
+from cutlass.cute.arch import clc as cute_clc
+
+# @@INJECT_TILE_CONSTANTS@@
+
+
+CLC_SCHED_STAGES = 2
+
+# Programmatic Dependent Launch (PDL, sm_90+).
+USE_PDL = True
+
+# Double-buffer for the TMA-store epilogue path
+EPI_SMEM_STAGES = 2
+
+# Named barrier id for cross-warp sync of the 4 epilogue warps
+EPI_SYNC_BAR_ID = 1
+
+# Named barrier id for the TMEM-alloc handoff
+TMEM_ALLOC_BARRIER_ID = 2
+
+
+@cute.jit
+def _auto_swizzle_w(m, n, k, nt_n):
+    """N-super-block width for the tile rasterization, resolved per launch.
+
+    ``tile_swizzle_n > 0`` pins it. Otherwise: the walk keeps one operand slice
+    resident and re-reads the other every super-block, so block along the SHORTER
+    problem side. Once that side outgrows what L2 can hold onto while C streams
+    through it, keeping it is no longer free -- fall back to the widest N block the
+    budget does cover.
+    """
+    if cutlass.const_expr(tile_swizzle_n > 0):
+        return tile_swizzle_n
+    budget = cutlass.Int64(swizzle_l2_budget_bytes)
+    row_bytes = (cutlass.Int64(ab_dtype.width) * k) // 8
+    cap = cutlass.max(budget // (row_bytes * cgrp_tile_mnk[1]), cutlass.Int64(1))
+    w = cutlass.min(cutlass.Int64(nt_n), cap)
+    if cutlass.min(m, n) * row_bytes <= budget and m <= n:
+        w = cutlass.Int64(1)
+    return cutlass.Int32(w)
+
+
+def _b_collector_op(mi):
+    """B is identical across the M sub-blocks (only A's address advances), so the
+    first MMA fills the B collector and the rest read it back instead of
+    re-fetching the same operand from SMEM."""
+    if cutlass.const_expr(not b_collector_ok or num_mma_m == 1):
+        return None
+    if cutlass.const_expr(mi == 0):
+        return nvvm.Tcgen05MMACollectorOp.FILL
+    if cutlass.const_expr(mi == num_mma_m - 1):
+        return nvvm.Tcgen05MMACollectorOp.LASTUSE
+    return nvvm.Tcgen05MMACollectorOp.USE
+
+
+@cute.kernel
+def _kernel(
+    m: cutlass.Int64,
+    n: cutlass.Int64,
+    k: cutlass.Int64,
+    # @@INJECT_KERNEL_AB_DESC_PARAMS@@
+    # @@INJECT_KERNEL_TAP_PARAMS@@
+    # @@INJECT_KERNEL_REDUCTION_STRIDE_PARAMS@@
+    # @@INJECT_KERNEL_AUX_PARAMS@@
+    # @@TMA_STORE_ONLY:BEGIN@@
+    # @@INJECT_KERNEL_TMA_C_PARAMS@@
+    # @@TMA_STORE_ONLY:END@@
+) -> None:
+    # @@INJECT_AB_DESC_LISTS@@
+    # @@TMA_STORE_ONLY:BEGIN@@
+    # @@INJECT_TMA_C_LISTS@@
+    tma_c_desc = tma_c_descs[0]
+    # @@TMA_STORE_ONLY:END@@
+
+    mma_warp_id = 4
+    tma_warp_id = 5
+    scheduler_warp_id = 6
+    unused_warp_id = 7
+    num_epilogue_warps = 4
+    epi_reg_count = 232
+    prod_reg_count = 24
+
+    warp_idx = cute.arch.warp_idx()
+    warp_idx = cute.arch.make_warp_uniform(warp_idx)
+    elect_one = nvvm.elect_sync()
+
+    tidx = cute.arch.thread_idx()[0]
+    bidx = cute.arch.block_idx()[0]
+    bidy = cute.arch.block_idx()[1]
+    bidz = cute.arch.block_idx()[2]
+    gridx = cute.arch.grid_dim()[0]
+    gridy = cute.arch.grid_dim()[1]
+
+    cluster_m = cluster_shape_mnk[0]
+    cluster_n = cluster_shape_mnk[1]
+    cluster_size = cluster_m * cluster_n * cluster_shape_mnk[2]
+
+    cta_rank_in_cluster = cute.arch.block_idx_in_cluster()
+    m_rank = cta_rank_in_cluster % cluster_m
+    n_rank = cta_rank_in_cluster // cluster_m
+
+    is_cluster_leader_cta = cta_rank_in_cluster == 0
+
+    full_cluster_mask = cutlass.Int16((1 << cluster_size) - 1)
+
+    if warp_idx == mma_warp_id:
+        for _i in cutlass.range_constexpr(num_a_operands):
+            nvvm.prefetch_tensormap(tma_a_descs[_i].get_ptr())
+            nvvm.prefetch_tensormap(tma_sfa_descs[_i].get_ptr())
+        for _j in cutlass.range_constexpr(num_b_operands):
+            nvvm.prefetch_tensormap(tma_b_descs[_j].get_ptr())
+            nvvm.prefetch_tensormap(tma_sfb_descs[_j].get_ptr())
+
+        # @@TMA_STORE_ONLY:BEGIN@@
+        nvvm.prefetch_tensormap(tma_c_desc.get_ptr())
+        # @@TMA_STORE_ONLY:END@@
+
+    swizzle_w = _auto_swizzle_w(m, n, k, gridy // cluster_n)
+    init_tile_m, init_tile_n = _l2_swizzle_tile(
+        bidx // cluster_m,
+        bidy // cluster_n,
+        gridx // cluster_m,
+        gridy // cluster_n,
+        swizzle_w,
+    )
+    init_tile_l = bidz
+
+    a_pattern = 0
+    for n_idx in cutlass.range_constexpr(cluster_n):
+        a_pattern = a_pattern | (1 << (n_idx * cluster_m))
+    b_pattern = (1 << cluster_m) - 1
+
+    if cutlass.const_expr(multicast_a):
+        tma_mcast_mask_a = cutlass.Int16(a_pattern) << m_rank
+    else:
+        tma_mcast_mask_a = cutlass.Int16(1) << cta_rank_in_cluster
+    if cutlass.const_expr(multicast_b):
+        tma_mcast_mask_b = cutlass.Int16(b_pattern) << (n_rank * cluster_m)
+    else:
+        tma_mcast_mask_b = cutlass.Int16(1) << cta_rank_in_cluster
+
+    a_part_arrive = cutlass.Int16(a_pattern) << m_rank
+    b_part_arrive = cutlass.Int16(b_pattern) << (n_rank * cluster_m)
+    ab_empty_arrive_mask = a_part_arrive | b_part_arrive
+
+    _smem_sys_reserved = cutlass.Array(cutlass.Int8, 1024, space=cutlass.AddressSpace.smem, alignment=1)
+
+    ab_full_mbar_ptr = cutlass.Array(cutlass.Int64, ab_stages, space=cutlass.AddressSpace.smem)
+    ab_empty_mbar_ptr = cutlass.Array(cutlass.Int64, ab_stages, space=cutlass.AddressSpace.smem)
+    acc_empty_mbar_ptr = cutlass.Array(cutlass.Int64, acc_stages, space=cutlass.AddressSpace.smem)
+    acc_full_mbar_ptr = cutlass.Array(cutlass.Int64, acc_stages, space=cutlass.AddressSpace.smem)
+    tmem_dealloc_mbar_ptr = cutlass.Array(cutlass.Int64, 1, space=cutlass.AddressSpace.smem)
+    tmem_ptr_i32 = cutlass.Array(cutlass.Int32, 1, space=cutlass.AddressSpace.smem)
+
+    _clc_response_raw = cutlass.Array(cutlass.Int128, CLC_SCHED_STAGES, space=cutlass.AddressSpace.smem, alignment=16)
+    clc_response_ptr_base = cute.make_ptr(
+        cutlass.Int128,
+        _clc_response_raw.data_ptr(),
+        mem_space=cute.AddressSpace.smem,
+    )
+    clc_full_mbar_ptr = cutlass.Array(cutlass.Int64, CLC_SCHED_STAGES, space=cutlass.AddressSpace.smem, alignment=8)
+    clc_empty_mbar_ptr = cutlass.Array(cutlass.Int64, CLC_SCHED_STAGES, space=cutlass.AddressSpace.smem, alignment=8)
+    clc_full_mbar_cute_base = cute.make_ptr(
+        cutlass.Int64,
+        clc_full_mbar_ptr.data_ptr(),
+        mem_space=cute.AddressSpace.smem,
+    )
+
+    sA_elems = sA_packed_elems
+    sB_elems = sB_packed_elems
+    smem_a_list = [
+        cutlass.Array(
+            ab_dtype,
+            sA_elems * ab_stages,
+            space=cutlass.AddressSpace.smem,
+            alignment=1024,
+        )
+        for _ in range(num_a_operands)
+    ]
+    smem_b_list = [
+        cutlass.Array(
+            ab_dtype,
+            sB_elems * ab_stages,
+            space=cutlass.AddressSpace.smem,
+            alignment=1024,
+        )
+        for _ in range(num_b_operands)
+    ]
+    smem_sfa_list = [
+        cutlass.Array(
+            cutlass.Uint8,
+            sfa_smem_bytes * ab_stages,
+            space=cutlass.AddressSpace.smem,
+            alignment=1024,
+        )
+        for _ in range(num_a_operands)
+    ]
+    smem_sfb_list = [
+        cutlass.Array(
+            cutlass.Uint8,
+            sfb_smem_bytes * ab_stages,
+            space=cutlass.AddressSpace.smem,
+            alignment=1024,
+        )
+        for _ in range(num_b_operands)
+    ]
+
+    # @@TMA_STORE_ONLY:BEGIN@@
+    epi_subtile_elems = epi_tile_mn[0] * epi_tile_mn[1]
+    smem_d_ptr = cutlass.Array(
+        cd_dtype,
+        epi_subtile_elems * EPI_SMEM_STAGES,
+        space=cutlass.AddressSpace.smem,
+        alignment=1024,
+    )
+    # @@TMA_STORE_ONLY:END@@
+
+    ab_empty_count = cluster_m + cluster_n - 1
+    num_consumer_warps_per_cta = 7
+    clc_empty_count = num_consumer_warps_per_cta * cluster_size
+    if warp_idx == 0:
+        for i in range(ab_stages):
+            if elect_one:
+                nvvm.mbarrier_init(ab_full_mbar_ptr.subview(i), 1)
+            if elect_one:
+                nvvm.mbarrier_init(ab_empty_mbar_ptr.subview(i), ab_empty_count)
+        for i in range(acc_stages):
+            if elect_one:
+                nvvm.mbarrier_init(acc_full_mbar_ptr.subview(i), 1)
+            if elect_one:
+                nvvm.mbarrier_init(acc_empty_mbar_ptr.subview(i), num_epilogue_warps)
+        if cutlass.const_expr(use_acc_overlap):
+            if elect_one:
+                nvvm.mbarrier_init(tmem_dealloc_mbar_ptr, num_epilogue_warps)
+        for i in range(CLC_SCHED_STAGES):
+            if elect_one:
+                nvvm.mbarrier_init(clc_full_mbar_ptr.subview(i), 1)
+            if elect_one:
+                nvvm.mbarrier_init(clc_empty_mbar_ptr.subview(i), clc_empty_count)
+    nvvm.fence_mbarrier_init()
+    if cutlass.const_expr(cluster_shape_mnk[0] * cluster_shape_mnk[1] > 1):
+        nvvm.barrier_cluster_arrive_relaxed()
+        nvvm.barrier_cluster_wait()
+    else:
+        nvvm.barrier_cta_sync(0)
+
+    sA_bytes = sA_elems * (ab_dtype.width // 8)
+    sB_bytes = sB_elems * (ab_dtype.width // 8)
+    num_tma_copy_bytes = num_a_operands * (sA_bytes + sfa_smem_bytes) + num_b_operands * (sB_bytes + sfb_smem_bytes)
+
+    cols_per_acc_stage = cta_tile_mnk[1]
+    tmem_alloc_bar_count = (num_epilogue_warps + 1) * 32
+
+    # @@INJECT_TAP_PTRS@@
+
+    VEC_BYTES = vec_bytes_epi
+    vsize = (VEC_BYTES * 8) // cd_dtype.width
+
+    M = m
+    N = n
+    num_tile_m = cute.ceil_div(M, cgrp_tile_mnk[0])
+    num_tile_n = cute.ceil_div(N, cgrp_tile_mnk[1])
+    num_k_tiles = cute.ceil_div(k, cta_tile_mnk[2])
+
+    if warp_idx == scheduler_warp_id:
+        nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
+        sched_iter = cutlass.Int32(0)
+        clc_empty_phase = cutlass.Int32(1)
+        clc_full_phase = cutlass.Int32(0)
+        is_valid_sched = cutlass.Int32(1)
+        while is_valid_sched != 0:
+            stage = sched_iter % CLC_SCHED_STAGES
+            if stage == 0 and sched_iter != 0:
+                clc_empty_phase = clc_empty_phase ^ 1
+                clc_full_phase = clc_full_phase ^ 1
+
+            if is_cluster_leader_cta:
+                while not nvvm.mbarrier_try_wait_parity(clc_empty_mbar_ptr.subview(stage), clc_empty_phase, time_limit=10_000_000):
+                    pass
+
+            if elect_one:
+                nvvm.mbarrier_arrive_expect_tx(clc_full_mbar_ptr.subview(stage), 16)
+
+            if is_cluster_leader_cta:
+                if elect_one:
+                    cute_clc.issue_clc_query(
+                        clc_full_mbar_cute_base + stage,
+                        clc_response_ptr_base + stage,
+                        multicast=True,
+                    )
+
+            while not nvvm.mbarrier_try_wait_parity(clc_full_mbar_ptr.subview(stage), clc_full_phase, time_limit=10_000_000):
+                pass
+
+            _m_idx, _n_idx, _l_idx, vld = cute_clc.clc_response(clc_response_ptr_base + stage)
+            cute.arch.fence_proxy("async.shared", space="cta")
+            is_valid_sched = vld
+
+            nvvm.bar_warp_sync(0xFFFFFFFF)
+            if elect_one:
+                empty_remote = nvvm.mapa(clc_empty_mbar_ptr.subview(stage), 0)
+                nvvm.mbarrier_arrive(empty_remote, scope=nvvm.MemScope.CLUSTER, relaxed=True)
+
+            sched_iter += 1
+
+        if cutlass.const_expr(cluster_shape_mnk[0] * cluster_shape_mnk[1] > 1):
+            if is_cluster_leader_cta:
+                for _ in range(CLC_SCHED_STAGES):
+                    stage = sched_iter % CLC_SCHED_STAGES
+                    if stage == 0 and sched_iter != 0:
+                        clc_empty_phase = clc_empty_phase ^ 1
+                    while not nvvm.mbarrier_try_wait_parity(
+                        clc_empty_mbar_ptr.subview(stage),
+                        clc_empty_phase,
+                        time_limit=10_000_000,
+                    ):
+                        pass
+                    sched_iter += 1
+
+    if warp_idx == tma_warp_id:
+        nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
+        if cutlass.const_expr(USE_PDL):
+            if elect_one:
+                nvvm.griddepcontrol("wait")
+        ab_empty_phase_bit = cutlass.Int32(1)
+        ab_iter = cutlass.Int32(0)
+        tile_m = init_tile_m
+        tile_n = init_tile_n
+        tile_l = init_tile_l
+        tile_iter = cutlass.Int32(0)
+        is_valid = cutlass.Int32(1)
+        clc_full_phase_tma = cutlass.Int32(0)
+        while is_valid != 0:
+            coord_m_per_cta = tile_m * cgrp_tile_mnk[0] + m_rank * cta_tile_mnk[0]
+            coord_n_per_cta = tile_n * cgrp_tile_mnk[1] + n_rank * cta_tile_mnk[1]
+            if cutlass.const_expr(matmul_a_batch == 1):
+                tile_l_a = cutlass.Int32(0)
+            else:
+                tile_l_a = tile_l
+            if cutlass.const_expr(matmul_b_batch == 1):
+                tile_l_b = cutlass.Int32(0)
+            else:
+                tile_l_b = tile_l
+
+            for k_tile_idx in range(num_k_tiles):
+                stage = ab_iter % ab_stages
+                if stage == 0 and ab_iter != 0:
+                    ab_empty_phase_bit = ab_empty_phase_bit ^ 1
+
+                while not nvvm.mbarrier_try_wait_parity(ab_empty_mbar_ptr.subview(stage), ab_empty_phase_bit, time_limit=10_000_000):
+                    pass
+
+                coord_k = k_tile_idx * cta_tile_mnk[2]
+                coord_sf_k = k_tile_idx * sf_tma_box_k
+                if elect_one:
+                    nvvm.mbarrier_arrive_expect_tx(ab_full_mbar_ptr.subview(stage), num_tma_copy_bytes)
+
+                for _ai in cutlass.range_constexpr(num_a_operands):
+                    sA_stage = smem_a_list[_ai].subview(sA_elems * stage)
+                    tma_a_desc = tma_a_descs[_ai]
+                    sSFA_stage = smem_sfa_list[_ai].subview(sfa_smem_bytes * stage)
+                    tma_sfa_desc = tma_sfa_descs[_ai]
+                    sfa_m_block = coord_m_per_cta // 128
+                    if cutlass.const_expr(multicast_a):
+                        if n_rank == 0:
+                            if cutlass.const_expr(a_is_m_major):
+                                for m_group in cutlass.range_constexpr(cta_tile_mnk[0] // a_tma_group_elems):
+                                    if elect_one:
+                                        nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                            sA_stage.subview(m_group * a_tma_group_elems * cta_tile_mnk[2]),
+                                            tma_a_desc.get_ptr(),
+                                            (
+                                                coord_m_per_cta + m_group * a_tma_group_elems,
+                                                coord_k,
+                                                tile_l_a,
+                                            ),
+                                            ab_full_mbar_ptr.subview(stage),
+                                            [],
+                                            multicast_mask=tma_mcast_mask_a,
+                                            group=nvvm.CTAGroup.CTA_1,
+                                        )
+                            else:
+                                if elect_one:
+                                    nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                        sA_stage,
+                                        tma_a_desc.get_ptr(),
+                                        (coord_k, coord_m_per_cta, tile_l_a),
+                                        ab_full_mbar_ptr.subview(stage),
+                                        [],
+                                        multicast_mask=tma_mcast_mask_a,
+                                        group=nvvm.CTAGroup.CTA_1,
+                                    )
+                    else:
+                        if cutlass.const_expr(a_is_m_major):
+                            for m_group in cutlass.range_constexpr(cta_tile_mnk[0] // a_tma_group_elems):
+                                if elect_one:
+                                    nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                        sA_stage.subview(m_group * a_tma_group_elems * cta_tile_mnk[2]),
+                                        tma_a_desc.get_ptr(),
+                                        (
+                                            coord_m_per_cta + m_group * a_tma_group_elems,
+                                            coord_k,
+                                            tile_l_a,
+                                        ),
+                                        ab_full_mbar_ptr.subview(stage),
+                                        [],
+                                        multicast_mask=tma_mcast_mask_a,
+                                        group=nvvm.CTAGroup.CTA_1,
+                                    )
+                        else:
+                            if elect_one:
+                                nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                    sA_stage,
+                                    tma_a_desc.get_ptr(),
+                                    (coord_k, coord_m_per_cta, tile_l_a),
+                                    ab_full_mbar_ptr.subview(stage),
+                                    [],
+                                    multicast_mask=tma_mcast_mask_a,
+                                    group=nvvm.CTAGroup.CTA_1,
+                                )
+                    if cutlass.const_expr(multicast_a):
+                        if n_rank == 0:
+                            if elect_one:
+                                nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                    sSFA_stage,
+                                    tma_sfa_desc.get_ptr(),
+                                    (0, coord_sf_k, sfa_m_block, tile_l_a),
+                                    ab_full_mbar_ptr.subview(stage),
+                                    [],
+                                    multicast_mask=tma_mcast_mask_a,
+                                    group=nvvm.CTAGroup.CTA_1,
+                                )
+                    else:
+                        if elect_one:
+                            nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                sSFA_stage,
+                                tma_sfa_desc.get_ptr(),
+                                (0, coord_sf_k, sfa_m_block, tile_l_a),
+                                ab_full_mbar_ptr.subview(stage),
+                                [],
+                                multicast_mask=tma_mcast_mask_a,
+                                group=nvvm.CTAGroup.CTA_1,
+                            )
+
+                for _bj in cutlass.range_constexpr(num_b_operands):
+                    sB_stage = smem_b_list[_bj].subview(sB_elems * stage)
+                    tma_b_desc = tma_b_descs[_bj]
+                    sSFB_stage = smem_sfb_list[_bj].subview(sfb_smem_bytes * stage)
+                    tma_sfb_desc = tma_sfb_descs[_bj]
+                    sfb_n_block = coord_n_per_cta // 128
+                    if cutlass.const_expr(multicast_b):
+                        if m_rank == 0:
+                            if cutlass.const_expr(b_is_n_major):
+                                for n_group in cutlass.range_constexpr(cta_tile_mnk[1] // b_tma_group_elems):
+                                    if elect_one:
+                                        nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                            sB_stage.subview(n_group * b_tma_group_elems * cta_tile_mnk[2]),
+                                            tma_b_desc.get_ptr(),
+                                            (
+                                                coord_n_per_cta + n_group * b_tma_group_elems,
+                                                coord_k,
+                                                tile_l_b,
+                                            ),
+                                            ab_full_mbar_ptr.subview(stage),
+                                            [],
+                                            multicast_mask=tma_mcast_mask_b,
+                                            group=nvvm.CTAGroup.CTA_1,
+                                        )
+                            else:
+                                if elect_one:
+                                    nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                        sB_stage,
+                                        tma_b_desc.get_ptr(),
+                                        (coord_k, coord_n_per_cta, tile_l_b),
+                                        ab_full_mbar_ptr.subview(stage),
+                                        [],
+                                        multicast_mask=tma_mcast_mask_b,
+                                        group=nvvm.CTAGroup.CTA_1,
+                                    )
+                    else:
+                        if cutlass.const_expr(b_is_n_major):
+                            for n_group in cutlass.range_constexpr(cta_tile_mnk[1] // b_tma_group_elems):
+                                if elect_one:
+                                    nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                        sB_stage.subview(n_group * b_tma_group_elems * cta_tile_mnk[2]),
+                                        tma_b_desc.get_ptr(),
+                                        (
+                                            coord_n_per_cta + n_group * b_tma_group_elems,
+                                            coord_k,
+                                            tile_l_b,
+                                        ),
+                                        ab_full_mbar_ptr.subview(stage),
+                                        [],
+                                        multicast_mask=tma_mcast_mask_b,
+                                        group=nvvm.CTAGroup.CTA_1,
+                                    )
+                        else:
+                            if elect_one:
+                                nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                    sB_stage,
+                                    tma_b_desc.get_ptr(),
+                                    (coord_k, coord_n_per_cta, tile_l_b),
+                                    ab_full_mbar_ptr.subview(stage),
+                                    [],
+                                    multicast_mask=tma_mcast_mask_b,
+                                    group=nvvm.CTAGroup.CTA_1,
+                                )
+                    if cutlass.const_expr(multicast_b):
+                        if m_rank == 0:
+                            if elect_one:
+                                nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                    sSFB_stage,
+                                    tma_sfb_desc.get_ptr(),
+                                    (0, coord_sf_k, sfb_n_block, tile_l_b),
+                                    ab_full_mbar_ptr.subview(stage),
+                                    [],
+                                    multicast_mask=tma_mcast_mask_b,
+                                    group=nvvm.CTAGroup.CTA_1,
+                                )
+                    else:
+                        if elect_one:
+                            nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                sSFB_stage,
+                                tma_sfb_desc.get_ptr(),
+                                (0, coord_sf_k, sfb_n_block, tile_l_b),
+                                ab_full_mbar_ptr.subview(stage),
+                                [],
+                                multicast_mask=tma_mcast_mask_b,
+                                group=nvvm.CTAGroup.CTA_1,
+                            )
+
+                ab_iter += 1
+
+            consumer_stage = tile_iter % CLC_SCHED_STAGES
+            if consumer_stage == 0 and tile_iter != 0:
+                clc_full_phase_tma = clc_full_phase_tma ^ 1
+            while not nvvm.mbarrier_try_wait_parity(
+                clc_full_mbar_ptr.subview(consumer_stage),
+                clc_full_phase_tma,
+                time_limit=10_000_000,
+            ):
+                pass
+            m_idx, n_idx, l_idx, vld = cute_clc.clc_response(clc_response_ptr_base + consumer_stage)
+            cute.arch.fence_proxy("async.shared", space="cta")
+            is_valid = vld
+            tile_m, tile_n = _l2_swizzle_tile(
+                m_idx // cluster_m,
+                n_idx // cluster_n,
+                gridx // cluster_m,
+                gridy // cluster_n,
+                swizzle_w,
+            )
+            tile_l = l_idx
+            nvvm.bar_warp_sync(0xFFFFFFFF)
+            if elect_one:
+                empty_remote = nvvm.mapa(clc_empty_mbar_ptr.subview(consumer_stage), 0)
+                nvvm.mbarrier_arrive(empty_remote, scope=nvvm.MemScope.CLUSTER, relaxed=True)
+            tile_iter += 1
+
+        tail_stage = ab_iter % ab_stages
+        tail_phase = ab_empty_phase_bit
+        if tail_stage == 0 and ab_iter != 0:
+            tail_phase = tail_phase ^ 1
+        for _ in range(ab_stages - 1):
+            tail_stage = tail_stage + 1
+            if tail_stage == ab_stages:
+                tail_stage = cutlass.Int32(0)
+                tail_phase = tail_phase ^ 1
+        if elect_one:
+            while not nvvm.mbarrier_try_wait_parity(ab_empty_mbar_ptr.subview(tail_stage), tail_phase, time_limit=10_000_000):
+                pass
+
+    if warp_idx == mma_warp_id:
+        nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
+        nvvm.tcgen05_alloc(
+            tmem_ptr_i32,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_1,
+        )
+        nvvm.bar_warp_sync(0xFFFFFFFF)
+        nvvm.barrier_cta_arrive(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
+        tmem_raw_addr = tmem_ptr_i32.load()
+        base_col_id_root = tmem_raw_addr & 0xFFFF
+        base_row_id = tmem_raw_addr >> 16
+        ab_full_phase_bit = cutlass.Int32(0)
+        ab_iter = cutlass.Int32(0)
+        acc_empty_phase_bit = cutlass.Int32(1)
+        tile_iter = cutlass.Int32(0)
+        is_valid = cutlass.Int32(1)
+        clc_full_phase_mma = cutlass.Int32(0)
+        acc_stage = cutlass.Int32(0)
+        # fp4 packs its K-mode into the OMMA descriptor's 2-bit split field; fp8
+        # keeps the MX descriptor's 1-bit one. Both are built once, outside the
+        # loops — the fields depend only on j (the scale id within a word).
+        if cutlass.const_expr(idesc_is_omma):
+            idesc_by_j = [
+                cutlass.experimental.primitives.Tcgen05MxOmmaInstrDesc.build(
+                    a_dtype=idesc_a_dtype,
+                    b_dtype=idesc_b_dtype,
+                    scale_format=sf_scale_format,
+                    n_dim=mma_n_dim,
+                    m_dim=mma_m_dim,
+                    a_major=mma_a_major,
+                    b_major=mma_b_major,
+                    a_sf_id=j * sf_scales_per_inst,
+                    b_sf_id=j * sf_scales_per_inst,
+                    k_dim=mma_k_dim_mode,
+                )
+                for j in range(sf_insts_per_atom)
+            ]
+        else:
+            idesc_by_j = [
+                cutlass.experimental.primitives.Tcgen05MxInstrDesc.build(
+                    a_dtype=idesc_a_dtype,
+                    b_dtype=idesc_b_dtype,
+                    scale_format=sf_scale_format,
+                    n_dim=mma_n_dim,
+                    m_dim=mma_m_dim,
+                    a_major=mma_a_major,
+                    b_major=mma_b_major,
+                    a_sf_id=j * sf_scales_per_inst,
+                    b_sf_id=j * sf_scales_per_inst,
+                    k_dim=mma_k_dim_mode,
+                )
+                for j in range(sf_insts_per_atom)
+            ]
+        s2t_shape, s2t_multicast = nvvm.S2TCopyMode.S2T_32x128b_WARPX4
+        sfa_tmem_bases = [(base_row_id << 16) | (base_col_id_root + sfa_col_bases[i]) for i in range(num_a_operands)]
+        sfb_tmem_bases = [(base_row_id << 16) | (base_col_id_root + sfb_col_bases[j]) for j in range(num_b_operands)]
+        sfa_scale_ptrs = [nvvm.make_tmem_ptr(b, cutlass.Float32) for b in sfa_tmem_bases]
+        sfb_scale_ptrs = [nvvm.make_tmem_ptr(b, cutlass.Float32) for b in sfb_tmem_bases]
+        # utccp destination per (MN-block, atom within the scale word). SFB is
+        # atom-MAJOR across the N-blocks because ONE instruction walks all of
+        # them; SFA is block-major because one instruction covers exactly one
+        # 128-row block, so that word has to be contiguous. Both collapse to
+        # the same addresses at a single block, and to sm100's layout at
+        # word_atoms == 1.
+        sfa_dst_ptrs = [
+            [
+                [nvvm.make_tmem_ptr(sfa_tmem_bases[i] + m * registers_per_block + a * registers_per_atom, cutlass.Float32) for a in range(word_atoms)]
+                for m in range(num_blocks_m)
+            ]
+            for i in range(num_a_operands)
+        ]
+        sfb_dst_ptrs = [
+            [
+                [nvvm.make_tmem_ptr(sfb_tmem_bases[j] + (a * num_blocks_n + m) * registers_per_atom, cutlass.Float32) for a in range(word_atoms)]
+                for m in range(num_blocks_n)
+            ]
+            for j in range(num_b_operands)
+        ]
+        while is_valid != 0:
+            acc_stage = tile_iter % acc_stages
+            if acc_stage == 0 and tile_iter != 0:
+                acc_empty_phase_bit = acc_empty_phase_bit ^ 1
+
+            while not nvvm.mbarrier_try_wait_parity(
+                acc_empty_mbar_ptr.subview(acc_stage),
+                acc_empty_phase_bit,
+                time_limit=10_000_000,
+            ):
+                pass
+
+            if cutlass.const_expr(use_acc_overlap):
+                acc_base_col = base_col_id_root + (tile_iter % 2) * acc_stage_stride
+            else:
+                acc_base_col = base_col_id_root + acc_stage * acc_region_cols
+            # One accumulator per (gemm, M block); M block mi sits
+            # epi_cols_per_mma_m columns further into its GEMM's region and reads
+            # SF word block mi (SF words are one per 128 rows).
+            acc_tmem_ptrs = [
+                [
+                    nvvm.make_tmem_ptr(
+                        (base_row_id << 16) | (acc_base_col + g * acc_gemm_stride + mi * epi_cols_per_mma_m),
+                        cutlass.Float32,
+                    )
+                    for mi in range(num_mma_m)
+                ]
+                for g in range(num_gemms)
+            ]
+
+            scale_d = cutlass.Boolean(False)
+            for k_tile_idx in range(num_k_tiles):
+                stage = ab_iter % ab_stages
+                if stage == 0 and ab_iter != 0:
+                    ab_full_phase_bit = ab_full_phase_bit ^ 1
+
+                while not nvvm.mbarrier_try_wait_parity(ab_full_mbar_ptr.subview(stage), ab_full_phase_bit, time_limit=10_000_000):
+                    pass
+
+                desc_a_bases = [
+                    cutlass.experimental.primitives.Tcgen05SmemDesc.build(
+                        start_address=smem_a_list[i].subview(sA_elems * stage),
+                        leading_byte_offset=a_smem_desc_leading_byte_offset,
+                        stride_byte_offset=a_smem_desc_stride_byte_offset,
+                        layout=ab_smem_swizzle,
+                    )
+                    for i in range(num_a_operands)
+                ]
+                desc_b_bases = [
+                    cutlass.experimental.primitives.Tcgen05SmemDesc.build(
+                        start_address=smem_b_list[j].subview(sB_elems * stage),
+                        leading_byte_offset=b_smem_desc_leading_byte_offset,
+                        stride_byte_offset=b_smem_desc_stride_byte_offset,
+                        layout=ab_smem_swizzle,
+                    )
+                    for j in range(num_b_operands)
+                ]
+                desc_sfa_bases = [
+                    cutlass.experimental.primitives.Tcgen05SmemDesc.build(
+                        start_address=smem_sfa_list[i].subview(sfa_smem_bytes * stage),
+                        leading_byte_offset=16,
+                        stride_byte_offset=128,
+                        layout=cutlass.experimental.primitives.Tcgen05SmemSwizzle.NONE,
+                    )
+                    for i in range(num_a_operands)
+                ]
+                desc_sfb_bases = [
+                    cutlass.experimental.primitives.Tcgen05SmemDesc.build(
+                        start_address=smem_sfb_list[j].subview(sfb_smem_bytes * stage),
+                        leading_byte_offset=16,
+                        stride_byte_offset=128,
+                        layout=cutlass.experimental.primitives.Tcgen05SmemSwizzle.NONE,
+                    )
+                    for j in range(num_b_operands)
+                ]
+
+                # One SF word per group of MMAs, refreshed right before they read
+                # it. A word spans word_atoms consecutive K-atoms in SMEM.
+                for atom_r in cutlass.range(num_sf_atoms, unroll_full=True):
+                    for _ai in cutlass.range_constexpr(num_a_operands):
+                        for _m in cutlass.range_constexpr(num_blocks_m):
+                            for _a in cutlass.range_constexpr(word_atoms):
+                                if elect_one:
+                                    nvvm.tcgen05_cp(
+                                        s2t_shape,
+                                        sfa_dst_ptrs[_ai][_m][_a],
+                                        desc_sfa_bases[_ai] + (sf_atom_desc_stride * (atom_r * word_atoms + _a) + sf_block_desc_stride * _m),
+                                        group=nvvm.CTAGroup.CTA_1,
+                                        multicast=s2t_multicast,
+                                    )
+                    for _bj in cutlass.range_constexpr(num_b_operands):
+                        for _m in cutlass.range_constexpr(num_blocks_n):
+                            for _a in cutlass.range_constexpr(word_atoms):
+                                if elect_one:
+                                    nvvm.tcgen05_cp(
+                                        s2t_shape,
+                                        sfb_dst_ptrs[_bj][_m][_a],
+                                        desc_sfb_bases[_bj] + (sf_atom_desc_stride * (atom_r * word_atoms + _a) + sf_block_desc_stride * _m),
+                                        group=nvvm.CTAGroup.CTA_1,
+                                        multicast=s2t_multicast,
+                                    )
+                    for j in cutlass.range_constexpr(sf_insts_per_atom):
+                        k_block_idx = atom_r * sf_insts_per_atom + j
+                        idesc_k = idesc_by_j[j]
+                        for g in cutlass.range_constexpr(num_gemms):
+                            _ai = gemm_a_idx[g]
+                            _bj = gemm_b_idx[g]
+                            desc_a_k = desc_a_bases[_ai].advance_start_address(a_smem_k_step_bytes * k_block_idx)
+                            desc_b = desc_b_bases[_bj].advance_start_address(b_smem_k_step_bytes * k_block_idx)
+                            for mi in cutlass.range_constexpr(num_mma_m):
+                                # The M sub-block offset is a whole SMEM swizzle atom, so the
+                                # descriptor's swizzle phase is preserved. B and its SF are
+                                # shared; A's SF word block follows the M block.
+                                desc_a = desc_a_k.advance_start_address(a_smem_m_step_bytes * mi)
+                                if elect_one:
+                                    nvvm.tcgen05_mma_block_scale(
+                                        mma_block_scale_kind,
+                                        nvvm.CTAGroup.CTA_1,
+                                        acc_tmem_ptrs[g][mi],
+                                        desc_a,
+                                        desc_b,
+                                        idesc_k,
+                                        enable_input_d=scale_d,
+                                        scale_a=sfa_dst_ptrs[_ai][mi][0],
+                                        scale_b=sfb_scale_ptrs[_bj],
+                                        scale_vec_size=scale_vec_size,
+                                        b_collector_op=_b_collector_op(mi),
+                                    )
+                        # Every accumulator sees scale_d=False on exactly the first
+                        # k_block of the tile, so the flip stays outside mi.
+                        scale_d = cutlass.Boolean(True)
+
+                if elect_one:
+                    nvvm.tcgen05_commit(
+                        ab_empty_mbar_ptr.subview(stage),
+                        multicast_mask=ab_empty_arrive_mask,
+                        group=nvvm.CTAGroup.CTA_1,
+                    )
+                ab_iter += 1
+
+            if elect_one:
+                nvvm.tcgen05_commit(
+                    acc_full_mbar_ptr.subview(acc_stage),
+                    group=nvvm.CTAGroup.CTA_1,
+                )
+
+            consumer_stage = tile_iter % CLC_SCHED_STAGES
+            if consumer_stage == 0 and tile_iter != 0:
+                clc_full_phase_mma = clc_full_phase_mma ^ 1
+            while not nvvm.mbarrier_try_wait_parity(
+                clc_full_mbar_ptr.subview(consumer_stage),
+                clc_full_phase_mma,
+                time_limit=10_000_000,
+            ):
+                pass
+            _m_idx, _n_idx, _l_idx, vld = cute_clc.clc_response(clc_response_ptr_base + consumer_stage)
+            cute.arch.fence_proxy("async.shared", space="cta")
+            is_valid = vld
+            nvvm.bar_warp_sync(0xFFFFFFFF)
+            if elect_one:
+                empty_remote = nvvm.mapa(clc_empty_mbar_ptr.subview(consumer_stage), 0)
+                nvvm.mbarrier_arrive(empty_remote, scope=nvvm.MemScope.CLUSTER, relaxed=True)
+            tile_iter += 1
+
+        if cutlass.const_expr(USE_PDL):
+            if elect_one:
+                nvvm.griddepcontrol("launch_dependents")
+        nvvm.tcgen05_relinquish_alloc_permit(group=nvvm.CTAGroup.CTA_1)
+        tail_stage = acc_stage
+        tail_phase = acc_empty_phase_bit
+        if elect_one:
+            for _ in range(acc_stages):
+                tail_stage = tail_stage + 1
+                if tail_stage == acc_stages:
+                    tail_stage = cutlass.Int32(0)
+                    tail_phase = tail_phase ^ 1
+                while not nvvm.mbarrier_try_wait_parity(acc_empty_mbar_ptr.subview(tail_stage), tail_phase, time_limit=10_000_000):
+                    pass
+            if cutlass.const_expr(use_acc_overlap):
+                while not nvvm.mbarrier_try_wait_parity(tmem_dealloc_mbar_ptr, 0, time_limit=10_000_000):
+                    pass
+
+        nvvm.bar_warp_sync(0xFFFFFFFF)
+        alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
+        nvvm.tcgen05_dealloc(
+            alloc_ptr,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_1,
+        )
+
+    if warp_idx < num_epilogue_warps:
+        nvvm.setmaxregister(epi_reg_count, nvvm.SetMaxRegisterAction.INCREASE)
+        nvvm.barrier_cta_sync(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
+        tmem_raw_addr = tmem_ptr_i32.load()
+        base_col_id_root = tmem_raw_addr & 0xFFFF
+        base_row_id = tmem_raw_addr >> 16
+        if cutlass.const_expr(USE_PDL):
+            nvvm.griddepcontrol("wait")
+        tile_iter = cutlass.Int32(0)
+        acc_full_phase_bit = cutlass.Int32(0)
+        tile_m = init_tile_m
+        tile_n = init_tile_n
+        tile_l = init_tile_l
+        is_valid = cutlass.Int32(1)
+        clc_full_phase_epi = cutlass.Int32(0)
+
+        if cutlass.const_expr(mma_inst_shape_mnk[0] == 64):
+            row_id_with_warp_offset = base_row_id
+        else:
+            row_id_with_warp_offset = base_row_id + warp_idx * 32
+
+        # One M block's accumulator columns are contiguous.
+        subtile_cnt = cute.ceil_div(epi_cols_per_mma_m, 32)
+        t2r_inst_repx = epi_tile_mn[1]
+        if cutlass.const_expr(mma_inst_shape_mnk[0] == 64):
+            shape = nvvm.Tcgen05LdStShape.SHAPE_16X32BX2
+            ld_half_off = 0
+        else:
+            shape = nvvm.Tcgen05LdStShape.SHAPE_32X32B
+            ld_half_off = None
+        lane = tidx % 32
+
+        # @@TMA_STORE_ONLY:BEGIN@@
+        epi_stage_idx = cutlass.Int32(EPI_SMEM_STAGES - 1)
+        # @@TMA_STORE_ONLY:END@@
+
+        while is_valid != 0:
+            coord_m_tile = tile_m * cgrp_tile_mnk[0] + m_rank * cta_tile_mnk[0]
+            coord_n = tile_n * cgrp_tile_mnk[1] + n_rank * cta_tile_mnk[1]
+
+            acc_stage = tile_iter % acc_stages
+            if acc_stage == 0 and tile_iter != 0:
+                acc_full_phase_bit = acc_full_phase_bit ^ 1
+
+            while not nvvm.mbarrier_try_wait_parity(acc_full_mbar_ptr.subview(acc_stage), acc_full_phase_bit, time_limit=10_000_000):
+                pass
+
+            if cutlass.const_expr(use_acc_overlap):
+                acc_buf_parity = tile_iter % 2
+                acc_base_col = base_col_id_root + acc_buf_parity * acc_stage_stride
+            else:
+                acc_buf_parity = cutlass.Int32(0)
+                acc_base_col = base_col_id_root + acc_stage * acc_region_cols
+            # One pass per MMA-M block over its own column region.
+            for mi in cutlass.range_constexpr(num_mma_m):
+                coord_m = coord_m_tile + mi * mma_inst_shape_mnk[0]
+                mi_col_base = acc_base_col + mi * epi_cols_per_mma_m
+                tmem_col_addr_gemms = [(row_id_with_warp_offset << 16) | (mi_col_base + g * acc_gemm_stride) for g in range(num_gemms)]
+
+                if cutlass.const_expr(mma_inst_shape_mnk[0] == 64):
+                    row = coord_m + warp_idx * 16 + lane
+                    row_active = lane < 16
+                else:
+                    row = coord_m + tidx
+                    row_active = True
+
+                # @@INJECT_AUX_VIEWS@@
+
+                for subtile_idx in cutlass.range(subtile_cnt, unroll_full=True):
+                    if cutlass.const_expr(use_acc_overlap):
+                        _sub = subtile_idx + (1 - acc_buf_parity) * (subtile_cnt - 1 - 2 * subtile_idx)
+                        subtile_col_offset = _sub * 32
+                    else:
+                        subtile_col_offset = subtile_idx * 32
+                    if cutlass.const_expr(not (use_tma_store_epi and cd_out_is_m_major)):
+                        c_rmem_vecs = []
+                        for g in cutlass.range_constexpr(num_gemms):
+                            tmem = cutlass.inttoptr(
+                                tmem_col_addr_gemms[g] + subtile_col_offset,
+                                6,
+                                cutlass.Float32,
+                            )
+                            c_rmem_vecs.append(nvvm.tcgen05_ld(shape, tmem, num=t2r_inst_repx, offset=ld_half_off))
+                        c_rmem_vec = c_rmem_vecs[0]
+
+                    if use_acc_overlap and (not cd_out_is_m_major) and mi == num_mma_m - 1 and subtile_idx == acc_overlap_subtiles - 1:
+                        nvvm.tcgen05_wait(kind=nvvm.Tcgen05Wait.LOAD)
+                        nvvm.tcgen05_fence(nvvm.Tcgen05Fence.BEFORE_THREAD_SYNC)
+                        if elect_one:
+                            nvvm.mbarrier_arrive(acc_empty_mbar_ptr.subview(acc_stage))
+
+                    col = coord_n + subtile_col_offset
+
+                    # @@TMA_STORE_ONLY:BEGIN@@
+                    epi_stage_idx = (epi_stage_idx + 1) % EPI_SMEM_STAGES
+                    smem_subtile_ptr = smem_d_ptr.subview(epi_stage_idx * epi_subtile_elems)
+                    smem_thr_ptr = smem_subtile_ptr.subview(tidx * t2r_inst_repx)
+
+                    if cutlass.const_expr(cd_out_is_m_major):
+                        ld_col = mi_col_base + subtile_col_offset
+                        for _h in cutlass.range(2, unroll_full=True):
+                            ld_row = base_row_id + warp_idx * 32 + _h * 16
+                            ld_addr = (ld_row << 16) | ld_col
+                            ld_tmem = cutlass.inttoptr(ld_addr, 6, cutlass.Float32)
+                            _lv = nvvm.tcgen05_ld(nvvm.Tcgen05LdStShape.SHAPE_16X256B, ld_tmem, num=4)
+                            vec_f32 = _lv
+                            col_j = col
+                            linear_idx = tile_l * out_stride_l_0 + row * out_stride_m_0 + col_j * out_stride_n_0
+
+                            # @@INJECT_EPILOGUE@@
+
+                            _i32 = vec_out.bitcast(cutlass.Int32)
+                            for _blk in cutlass.range_constexpr(2):
+                                _regs = [_i32[_blk * 4 + _j] for _j in range(4)]
+                                _n_full = (lane % 8) + 8 * (lane // 16) + 16 * _blk
+                                _m_base = warp_idx * 32 + _h * 16 + 8 * ((lane // 8) % 2)
+                                _stm_off = (
+                                    (_m_base // cd_mmajor_atom_m) * (cd_mmajor_atom_m * epi_tile_mn[1])
+                                    + (_m_base % cd_mmajor_atom_m)
+                                    + _n_full * cd_mmajor_atom_m
+                                )
+                                nvvm.stmatrix(
+                                    _apply_smem_swizzle(
+                                        smem_subtile_ptr.data_ptr() + _stm_off,
+                                        cutlass.Swizzle(3, 4, 3),
+                                    ),
+                                    _regs,
+                                    nvvm.MMALayout.COL,
+                                    shape=nvvm.StoreShape.M8N8,
+                                )
+                    else:
+                        vec_f32 = c_rmem_vec
+                        col_j = col
+                        linear_idx = tile_l * out_stride_l_0 + row * out_stride_m_0 + col_j * out_stride_n_0
+
+                        # @@INJECT_EPILOGUE@@
+
+                        smem_thr_ptr.data_ptr().store_swizzled(vec_out, alignment=64, swizzle=cutlass.Swizzle(2, 4, 3))
+
+                    cute.arch.fence_view_async_shared()
+                    nvvm.barrier_cta_sync(
+                        barrier_id=EPI_SYNC_BAR_ID,
+                        thread_count=num_epilogue_warps * 32,
+                    )
+
+                    if warp_idx == 0:
+                        if cutlass.const_expr(cd_out_is_m_major):
+                            for _mb in cutlass.range_constexpr(epi_tile_mn[0] // cd_mmajor_atom_m):
+                                if elect_one:
+                                    nvvm.cp_async_bulk_tensor_global_shared_cta(
+                                        tma_c_desc.get_ptr(),
+                                        smem_subtile_ptr.subview(_mb * (cd_mmajor_atom_m * epi_tile_mn[1])),
+                                        (coord_m + _mb * cd_mmajor_atom_m, col, tile_l),
+                                    )
+                        else:
+                            if elect_one:
+                                nvvm.cp_async_bulk_tensor_global_shared_cta(
+                                    tma_c_desc.get_ptr(),
+                                    smem_subtile_ptr,
+                                    (col, coord_m, tile_l),
+                                )
+                        if elect_one:
+                            nvvm.cp_async_bulk_commit_group()
+                        nvvm.cp_async_bulk_wait_group(EPI_SMEM_STAGES - 1, read=True)
+
+                    nvvm.barrier_cta_sync(
+                        barrier_id=EPI_SYNC_BAR_ID,
+                        thread_count=num_epilogue_warps * 32,
+                    )
+                    # @@TMA_STORE_ONLY:END@@
+
+                    # @@STG_ONLY:BEGIN@@
+                    if row_active and row < M:
+                        for j in cutlass.range_constexpr(t2r_inst_repx // vsize):
+                            col_j = col + j * vsize
+                            if col_j + vsize <= N:
+                                vec_f32 = c_rmem_vec[j * vsize : (j + 1) * vsize]
+
+                                # @@INJECT_STG_VEC_BINDINGS@@
+
+                                # @@INJECT_EPILOGUE@@
+                    # @@STG_ONLY:END@@
+
+            if cutlass.const_expr((not use_acc_overlap) or cd_out_is_m_major):
+                nvvm.tcgen05_wait(kind=nvvm.Tcgen05Wait.LOAD)
+                nvvm.tcgen05_fence(nvvm.Tcgen05Fence.BEFORE_THREAD_SYNC)
+                if elect_one:
+                    nvvm.mbarrier_arrive(acc_empty_mbar_ptr.subview(acc_stage))
+
+            consumer_stage = tile_iter % CLC_SCHED_STAGES
+            if consumer_stage == 0 and tile_iter != 0:
+                clc_full_phase_epi = clc_full_phase_epi ^ 1
+            while not nvvm.mbarrier_try_wait_parity(
+                clc_full_mbar_ptr.subview(consumer_stage),
+                clc_full_phase_epi,
+                time_limit=10_000_000,
+            ):
+                pass
+            m_idx, n_idx, l_idx, vld = cute_clc.clc_response(clc_response_ptr_base + consumer_stage)
+            cute.arch.fence_proxy("async.shared", space="cta")
+            is_valid = vld
+            tile_m, tile_n = _l2_swizzle_tile(
+                m_idx // cluster_m,
+                n_idx // cluster_n,
+                gridx // cluster_m,
+                gridy // cluster_n,
+                swizzle_w,
+            )
+            tile_l = l_idx
+            nvvm.bar_warp_sync(0xFFFFFFFF)
+            if elect_one:
+                empty_remote = nvvm.mapa(clc_empty_mbar_ptr.subview(consumer_stage), 0)
+                nvvm.mbarrier_arrive(empty_remote, scope=nvvm.MemScope.CLUSTER, relaxed=True)
+
+            tile_iter += 1
+
+        if cutlass.const_expr(use_acc_overlap):
+            nvvm.tcgen05_wait(kind=nvvm.Tcgen05Wait.LOAD)
+            nvvm.tcgen05_fence(nvvm.Tcgen05Fence.BEFORE_THREAD_SYNC)
+            if elect_one:
+                nvvm.mbarrier_arrive(tmem_dealloc_mbar_ptr)
+
+        # @@TMA_STORE_ONLY:BEGIN@@
+        if warp_idx == 0:
+            nvvm.cp_async_bulk_wait_group(0, read=True)
+        # @@TMA_STORE_ONLY:END@@
+
+    if warp_idx == unused_warp_id:
+        nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
+
+
+@cute.jit
+def _host(
+    problem_size: tuple,
+    # @@INJECT_HOST_AB_PARAMS@@
+    # @@INJECT_HOST_TAP_PARAMS@@
+    # @@INJECT_HOST_AUX_PARAMS@@
+    # @@TMA_STORE_ONLY:BEGIN@@
+    # @@INJECT_HOST_TMA_C_PARAMS@@
+    # @@TMA_STORE_ONLY:END@@
+    stream: _cuda.CUstream,
+) -> None:
+    # @@INJECT_HOST_AB_LISTS@@
+
+    m = problem_size[0]
+    n = problem_size[1]
+    k_sym = problem_size[2]
+    batch = problem_size[3]
+    a_stride_m = problem_size[4]
+    a_stride_k = problem_size[5]
+    a_stride_l = problem_size[6]
+    b_stride_n = problem_size[7]
+    b_stride_k = problem_size[8]
+    b_stride_l = problem_size[9]
+
+    # @@INJECT_HOST_REDUCTION_STRIDES@@
+
+    if cutlass.const_expr(matmul_a_batch == 1):
+        a_batch = 1
+    else:
+        a_batch = batch
+    if cutlass.const_expr(matmul_b_batch == 1):
+        b_batch = 1
+    else:
+        b_batch = batch
+    rest_k = ((k_sym // block_size) + 3) // 4
+    rest_m = (m + 127) // 128
+    rest_n = (n + 127) // 128
+    tma_a_desc_list = []
+    tma_sfa_desc_list = []
+    for _a_op, _sfa_op in zip(_a_operands, _sfa_operands):
+        if cutlass.const_expr(a_is_m_major):
+            tma_a_desc_list.append(
+                _tma.create_tensor_map_tiled(
+                    global_address=_a_op.iterator.toint(),
+                    dtype=ab_tma_desc_dtype,
+                    global_dims=[m, k_sym, a_batch],
+                    global_strides=[
+                        a_stride_k * ab_dtype.width // 128,
+                        a_stride_l * ab_dtype.width // 128,
+                    ],
+                    box_dims=[a_tma_group_elems, cta_tile_mnk[2], 1],
+                    swizzle=ab_tma_swizzle,
+                    tma_format=ab_tma_format,
+                )
+            )
+        else:
+            tma_a_desc_list.append(
+                _tma.create_tensor_map_tiled(
+                    global_address=_a_op.iterator.toint(),
+                    dtype=ab_tma_desc_dtype,
+                    global_dims=[k_sym, m, a_batch],
+                    global_strides=[
+                        a_stride_m * ab_dtype.width // 128,
+                        a_stride_l * ab_dtype.width // 128,
+                    ],
+                    box_dims=[cta_tile_mnk[2], cta_tile_mnk[0], 1],
+                    swizzle=ab_tma_swizzle,
+                    tma_format=ab_tma_format,
+                )
+            )
+        sfa_fp16_tensor = cute.make_tensor(
+            cute.recast_ptr(_sfa_op.iterator, dtype=cutlass.Float16),
+            cute.make_layout(
+                (256, rest_k, rest_m, batch),
+                stride=(
+                    1,
+                    256,
+                    cute.assume(256 * rest_k, 8),
+                    cute.assume(256 * rest_k * rest_m, 8),
+                ),
+            ),
+        )
+        tma_sfa_desc_list.append(
+            _tma.create_tensor_map_tiled_from_view(
+                sfa_fp16_tensor,
+                dtype=cutlass.Uint16,
+                box_dims=(256, sf_tma_box_k, sfa_tma_box_mn, 1),
+                stride_order=(0, 1, 2, 3),
+                swizzle=_tma.TensorMapSwizzle.none,
+            )
+        )
+    tma_b_desc_list = []
+    tma_sfb_desc_list = []
+    for _b_op, _sfb_op in zip(_b_operands, _sfb_operands):
+        if cutlass.const_expr(b_is_n_major):
+            tma_b_desc_list.append(
+                _tma.create_tensor_map_tiled(
+                    global_address=_b_op.iterator.toint(),
+                    dtype=ab_tma_desc_dtype,
+                    global_dims=[n, k_sym, b_batch],
+                    global_strides=[
+                        b_stride_k * ab_dtype.width // 128,
+                        b_stride_l * ab_dtype.width // 128,
+                    ],
+                    box_dims=[b_tma_group_elems, cta_tile_mnk[2], 1],
+                    swizzle=ab_tma_swizzle,
+                    tma_format=ab_tma_format,
+                )
+            )
+        else:
+            tma_b_desc_list.append(
+                _tma.create_tensor_map_tiled(
+                    global_address=_b_op.iterator.toint(),
+                    dtype=ab_tma_desc_dtype,
+                    global_dims=[k_sym, n, b_batch],
+                    global_strides=[
+                        b_stride_n * ab_dtype.width // 128,
+                        b_stride_l * ab_dtype.width // 128,
+                    ],
+                    box_dims=[cta_tile_mnk[2], cta_tile_mnk[1], 1],
+                    swizzle=ab_tma_swizzle,
+                    tma_format=ab_tma_format,
+                )
+            )
+        sfb_fp16_tensor = cute.make_tensor(
+            cute.recast_ptr(_sfb_op.iterator, dtype=cutlass.Float16),
+            cute.make_layout(
+                (256, rest_k, rest_n, batch),
+                stride=(
+                    1,
+                    256,
+                    cute.assume(256 * rest_k, 8),
+                    cute.assume(256 * rest_k * rest_n, 8),
+                ),
+            ),
+        )
+        tma_sfb_desc_list.append(
+            _tma.create_tensor_map_tiled_from_view(
+                sfb_fp16_tensor,
+                dtype=cutlass.Uint16,
+                box_dims=(256, sf_tma_box_k, sfb_tma_box_mn, 1),
+                stride_order=(0, 1, 2, 3),
+                swizzle=_tma.TensorMapSwizzle.none,
+            )
+        )
+    # @@TMA_STORE_ONLY:BEGIN@@
+    # @@INJECT_HOST_TMA_C_LISTS@@
+    c = _tma_c_outputs[0]
+    if cutlass.const_expr(cd_out_is_m_major):
+        tma_c_desc = _tma.create_tensor_map_tiled(
+            global_address=c.iterator.toint(),
+            dtype=cd_tma_dtype,
+            global_dims=[m, n, batch],
+            global_strides=[
+                out_stride_n_0 * cd_dtype.width // 128,
+                out_stride_l_0 * cd_dtype.width // 128,
+            ],
+            box_dims=[cd_mmajor_atom_m, epi_tile_mn[1], 1],
+            swizzle=(_tma.TensorMapSwizzle.s128b if cutlass.const_expr(use_tma_store_epi) else _tma.TensorMapSwizzle.none),
+        )
+    else:
+        tma_c_desc = _tma.create_tensor_map_tiled(
+            global_address=c.iterator.toint(),
+            dtype=cd_tma_dtype,
+            global_dims=[n, m, batch],
+            global_strides=[
+                out_stride_m_0 * cd_dtype.width // 128,
+                out_stride_l_0 * cd_dtype.width // 128,
+            ],
+            box_dims=[epi_tile_mn[1], epi_tile_mn[0], 1],
+            swizzle=(_tma.TensorMapSwizzle.s64b if cutlass.const_expr(use_tma_store_epi) else _tma.TensorMapSwizzle.none),
+        )
+    tma_c_desc_list = [tma_c_desc]
+    # @@TMA_STORE_ONLY:END@@
+
+    cluster_m = cluster_shape_mnk[0]
+    cluster_n = cluster_shape_mnk[1]
+    cgrp_tile_m = cgrp_tile_mnk[0]
+    cgrp_tile_n = cgrp_tile_mnk[1]
+    num_tile_m_host = (m + cgrp_tile_m - 1) // cgrp_tile_m
+    num_tile_n_host = (n + cgrp_tile_n - 1) // cgrp_tile_n
+    grid_x = num_tile_m_host * cluster_m
+    grid_y = num_tile_n_host * cluster_n
+    grid_shape = (grid_x, grid_y, batch)
+    _kernel(
+        problem_size[0],
+        problem_size[1],
+        problem_size[2],
+        # @@INJECT_HOST_KERNEL_DESC_PASS@@
+        # @@INJECT_HOST_TAP_PASS@@
+        # @@INJECT_HOST_REDUCTION_STRIDE_PASS@@
+        # @@INJECT_HOST_AUX_PASS@@
+        # @@TMA_STORE_ONLY:BEGIN@@
+        # @@INJECT_HOST_TMA_C_PASS@@
+        # @@TMA_STORE_ONLY:END@@
+    ).launch(
+        grid=grid_shape,
+        block=(threads_per_cta, 1, 1),
+        cluster=cluster_shape_mnk,
+        use_pdl=USE_PDL,
+        stream=stream,
+    )
+
+
+@lru_cache(maxsize=None)
+def compile() -> Callable:
+    out_vec_elems = vec_bytes_epi // (cd_dtype.width // 8)
+    ab_stride_elems = 128 // ab_dtype.width
+    sym_m = cute.sym_int64()
+    sym_n = cute.sym_int64(divisibility=out_vec_elems)
+    # K tails are supported: the K loop is ceil_div and the TMA descriptor's global K
+    # extent makes a partial box HW zero-filled. The only real K rule is the 16-byte
+    # TMA contiguous-extent one, already gated by _tma_alignment_reject.
+    sym_k = cute.sym_int64()
+    # Packed K extent: same reasoning as sym_k -- no CTA-tile multiple is required.
+    sym_kp = cute.sym_int64()
+    sym_l = cute.sym_int64()
+    if matmul_a_batch == 1:
+        sym_a_l = 1
+    else:
+        sym_a_l = sym_l
+    if matmul_b_batch == 1:
+        sym_b_l = 1
+    else:
+        sym_b_l = sym_l
+
+    def _make_fake_a():
+        return make_fake_compact_tensor(
+            a_fake_dtype,
+            (sym_m, sym_kp, sym_a_l),
+            stride_order=(0, 1, 2) if a_is_m_major else (1, 0, 2),
+            assumed_align=16,
+        )
+
+    def _make_fake_b():
+        return make_fake_compact_tensor(
+            b_fake_dtype,
+            (sym_n, sym_kp, sym_b_l),
+            stride_order=(0, 1, 2) if b_is_n_major else (1, 0, 2),
+            assumed_align=16,
+        )
+
+    # SF reaches the kernel as a base pointer only; the host rebuilds the
+    # F8_128x4 view from problem_size, so no SF mode carries a layout contract.
+    def _make_fake_sfa():
+        return cute.runtime.make_fake_tensor(
+            sf_cutlass_dtype,
+            (cute.sym_int64(), cute.sym_int64(), cute.sym_int64()),
+            stride=(cute.sym_int64(), cute.sym_int64(), cute.sym_int64()),
+            assumed_align=16,
+        )
+
+    def _make_fake_sfb():
+        return cute.runtime.make_fake_tensor(
+            sf_cutlass_dtype,
+            (cute.sym_int64(), cute.sym_int64(), cute.sym_int64()),
+            stride=(cute.sym_int64(), cute.sym_int64(), cute.sym_int64()),
+            assumed_align=16,
+        )
+
+    # @@TMA_STORE_ONLY:BEGIN@@
+    def _make_fake_c():
+        return make_fake_compact_tensor(
+            cd_dtype,
+            (sym_m, sym_n // cd_fake_n_div, sym_l),
+            stride_order=(0, 1, 2) if cd_out_is_m_major else (1, 0, 2),
+            assumed_align=16,
+        )
+
+    # @@INJECT_COMPILE_TMA_C_FAKES@@
+    # @@TMA_STORE_ONLY:END@@
+
+    # @@INJECT_COMPILE_AB_FAKES@@
+
+    # The operand's unit stride (m/n when MN-major, k when K-major) never reaches TMA, so it carries no 16B contract.
+    sym_a_stride_m = cute.sym_int64() if a_is_m_major else cute.sym_int64(divisibility=ab_stride_elems)
+    sym_a_stride_k = cute.sym_int64(divisibility=ab_stride_elems) if a_is_m_major else cute.sym_int64()
+    sym_a_stride_l = cute.sym_int64(divisibility=ab_stride_elems)
+    sym_b_stride_n = cute.sym_int64() if b_is_n_major else cute.sym_int64(divisibility=ab_stride_elems)
+    sym_b_stride_k = cute.sym_int64(divisibility=ab_stride_elems) if b_is_n_major else cute.sym_int64()
+    sym_b_stride_l = cute.sym_int64(divisibility=ab_stride_elems)
+
+    # @@INJECT_COMPILE_REDUCTION_STRIDE_DECLS@@
+
+    # @@INJECT_COMPILE_TAP_FAKES@@
+
+    problem_size = (
+        sym_m,
+        sym_n,
+        sym_k,
+        sym_l,
+        sym_a_stride_m,
+        sym_a_stride_k,
+        sym_a_stride_l,
+        sym_b_stride_n,
+        sym_b_stride_k,
+        sym_b_stride_l,
+        # @@INJECT_COMPILE_REDUCTION_STRIDE_SYMBOLS@@
+    )
+
+    # @@INJECT_COMPILE_AUX_FAKES@@
+
+    _fake_stream = make_fake_stream(use_tvm_ffi_env_stream=False)
+    return cute.compile(
+        _host,
+        problem_size,
+        # @@INJECT_COMPILE_AB_PASS@@
+        # @@INJECT_COMPILE_TAP_PASS@@
+        # @@INJECT_COMPILE_AUX_PASS@@
+        # @@TMA_STORE_ONLY:BEGIN@@
+        # @@INJECT_COMPILE_TMA_C_PASS@@
+        # @@TMA_STORE_ONLY:END@@
+        stream=_fake_stream,
+        options=frost_compile_options,
+    )
diff --git a/python/cudnn/gemm/frost/kernel_templates/sm107_block_scale_matmul_2ctamma.py b/python/cudnn/gemm/frost/kernel_templates/sm107_block_scale_matmul_2ctamma.py
new file mode 100644
index 000000000..873f5c19c
--- /dev/null
+++ b/python/cudnn/gemm/frost/kernel_templates/sm107_block_scale_matmul_2ctamma.py
@@ -0,0 +1,1522 @@
+# Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: MIT
+
+"""sm107 cta_group=2 **block-scaled** GEMM kernel: cluster-pair MMA + CLC dynamic scheduler (2-stage ring) + double-TMEM.
+
+Computes ``C = (descale_a ⊙ A) @ (descale_b ⊙ B)`` with per-block scale factors
+applied inside the MMA (``tcgen05.mma...block_scale``); supports nvfp4 / mxfp4 /
+mxfp8. 2-CTA MMA pair (leader dispatches, follower CLC-consumes). Compiler picks
+this when ``TileConfig.cta_group == 2``.
+
+The pipeline is the sm100 one; SM 10.7's block-scale MMA reads a **64-byte K**
+per instruction instead of 32, which shows up in exactly two places (both
+driven by injected constants, so the rest of the file stays in lockstep with
+``sm100_block_scale_matmul_2ctamma.py``):
+
+  * half as many MMAs per K-tile, each consuming ``sf_scales_per_inst`` scales
+    (8 at K-block 16, 4 at 32 — it follows the BLOCK SIZE, not the scale
+    dtype). When that exceeds the 4 scales one 128x4 utccp atom holds, a scale
+    *word* spans ``word_atoms`` atoms, and the two SF regions then lay them
+    out DIFFERENTLY: SFB atom-major across its N-blocks, SFA block-major.
+    At K-block 32 ``word_atoms == 1`` (identical to sm100).
+  * fp4 rides the OMMA instruction descriptor (``Tcgen05MxOmmaInstrDesc``,
+    K-mode 2 = 128 fp4 elements); mxfp8 stays on ``Tcgen05MxInstrDesc``
+    (K-mode 1 = 64 fp8 elements). Both take the real operand dtype.
+
+(The 576-column TMEM this GPU has is an ARCH property, not a pipeline one —
+`_TMEM_COLS_BY_ARCH` hands it to the sm100 templates on the same part too.)
+
+Warp layout (8 warps × 32 = 256 threads/CTA):
+  warps 0–3 : epilogue (warp 0 also allocates TMEM)  — setmaxnreg.inc 216
+  warp  4   : MMA driver (leader CTA runs MMA; follower CTA CLC-consumes only)  — setmaxnreg.dec 40
+  warp  5   : TMA producer (both CTAs load their slice)  — setmaxnreg.dec 40
+  warp  6   : CLC scheduler (leader CTA issues queries; every CTA waits + reads + arrives empty)  — setmaxnreg.dec 40
+  warp  7   : unused donor  — setmaxnreg.dec 40
+"""
+
+from __future__ import annotations
+
+from functools import lru_cache
+from typing import Callable
+
+import cutlass.experimental.primitives as nvvm
+from cudnn.gemm.frost.kernel_templates._tile_helpers import (
+    l2_swizzle_tile as _l2_swizzle_tile,
+)
+import cutlass.experimental.cuda.tensor_map as _tma
+import cutlass._mlir_helpers.vector as _cvec
+from cutlass import apply_swizzle as _apply_smem_swizzle
+import cutlass
+import cutlass.cute as cute
+from cutlass.cute.runtime import make_fake_compact_tensor
+from cutlass.cute.runtime import make_fake_stream
+from cuda.bindings import driver as _cuda
+from cutlass.cute.arch import clc as cute_clc
+
+# @@INJECT_TILE_CONSTANTS@@
+
+
+# Scheduler ring depth.
+CLC_SCHED_STAGES = 2
+
+# Programmatic Dependent Launch (PDL, sm_90+).
+USE_PDL = True
+
+# Double-buffer for the TMA-store epilogue path
+EPI_SMEM_STAGES = 2
+
+# Named barrier id for the 4-warp epilogue handoff around the TMA store.
+EPI_SYNC_BAR_ID = 1
+
+# Named barrier id for the TMEM-alloc handoff
+TMEM_ALLOC_BARRIER_ID = 2
+
+
+@cute.jit
+def _auto_swizzle_w(m, n, k, nt_n):
+    """N-super-block width for the tile rasterization, resolved per launch.
+
+    ``tile_swizzle_n > 0`` pins it. Otherwise: the walk keeps one operand slice
+    resident and re-reads the other every super-block, so block along the SHORTER
+    problem side. Once that side outgrows what L2 can hold onto while C streams
+    through it, keeping it is no longer free -- fall back to the widest N block the
+    budget does cover.
+    """
+    if cutlass.const_expr(tile_swizzle_n > 0):
+        return tile_swizzle_n
+    budget = cutlass.Int64(swizzle_l2_budget_bytes)
+    row_bytes = (cutlass.Int64(ab_dtype.width) * k) // 8
+    cap = cutlass.max(budget // (row_bytes * cgrp_tile_mnk[1]), cutlass.Int64(1))
+    w = cutlass.min(cutlass.Int64(nt_n), cap)
+    if cutlass.min(m, n) * row_bytes <= budget and m <= n:
+        w = cutlass.Int64(1)
+    return cutlass.Int32(w)
+
+
+def _b_collector_op(mi):
+    """B is identical across the M sub-blocks (only A's address advances), so the
+    first MMA fills the B collector and the rest read it back instead of
+    re-fetching the same operand from SMEM."""
+    if cutlass.const_expr(not b_collector_ok or num_mma_m == 1):
+        return None
+    if cutlass.const_expr(mi == 0):
+        return nvvm.Tcgen05MMACollectorOp.FILL
+    if cutlass.const_expr(mi == num_mma_m - 1):
+        return nvvm.Tcgen05MMACollectorOp.LASTUSE
+    return nvvm.Tcgen05MMACollectorOp.USE
+
+
+@cute.kernel
+def _kernel(
+    m: cutlass.Int64,
+    n: cutlass.Int64,
+    k: cutlass.Int64,
+    # @@INJECT_KERNEL_AB_DESC_PARAMS@@
+    # @@INJECT_KERNEL_TAP_PARAMS@@
+    # @@INJECT_KERNEL_REDUCTION_STRIDE_PARAMS@@
+    # @@INJECT_KERNEL_AUX_PARAMS@@
+    # @@TMA_STORE_ONLY:BEGIN@@
+    # @@INJECT_KERNEL_TMA_C_PARAMS@@
+    # @@TMA_STORE_ONLY:END@@
+) -> None:
+    # @@INJECT_AB_DESC_LISTS@@
+    # @@TMA_STORE_ONLY:BEGIN@@
+    # @@INJECT_TMA_C_LISTS@@
+    tma_c_desc = tma_c_descs[0]
+    # @@TMA_STORE_ONLY:END@@
+
+    mma_warp_id = 4
+    tma_warp_id = 5
+    scheduler_warp_id = 6
+    unused_warp_id = 7
+    num_epilogue_warps = 4
+    epi_reg_count = 232
+    prod_reg_count = 24
+
+    warp_idx = cute.arch.warp_idx()
+    warp_idx = cute.arch.make_warp_uniform(warp_idx)
+    elect_one = nvvm.elect_sync()
+
+    tidx = cute.arch.thread_idx()[0]
+    bidx = cute.arch.block_idx()[0]
+    bidy = cute.arch.block_idx()[1]
+    bidz = cute.arch.block_idx()[2]
+    gridx = cute.arch.grid_dim()[0]
+    gridy = cute.arch.grid_dim()[1]
+
+    cluster_m = cluster_shape_mnk[0]
+    cluster_n = cluster_shape_mnk[1]
+    cluster_size = cluster_m * cluster_n * cluster_shape_mnk[2]
+
+    cta_rank_in_cluster = cute.arch.block_idx_in_cluster()
+    m_rank = cta_rank_in_cluster % cluster_m
+    n_rank = cta_rank_in_cluster // cluster_m
+    pair_member = m_rank % 2
+    pair_m_idx = m_rank // 2
+    is_pair_leader = pair_member == 0
+    pair_leader_rank = pair_m_idx * 2 + n_rank * cluster_m
+
+    is_cluster_leader_cta = cta_rank_in_cluster == 0
+
+    if warp_idx == mma_warp_id:
+        for _i in cutlass.range_constexpr(num_a_operands):
+            nvvm.prefetch_tensormap(tma_a_descs[_i].get_ptr())
+            nvvm.prefetch_tensormap(tma_sfa_descs[_i].get_ptr())
+        for _j in cutlass.range_constexpr(num_b_operands):
+            nvvm.prefetch_tensormap(tma_b_descs[_j].get_ptr())
+            nvvm.prefetch_tensormap(tma_sfb_descs[_j].get_ptr())
+
+        # @@TMA_STORE_ONLY:BEGIN@@
+        nvvm.prefetch_tensormap(tma_c_desc.get_ptr())
+        # @@TMA_STORE_ONLY:END@@
+
+    swizzle_w = _auto_swizzle_w(m, n, k, gridy // cluster_n)
+    init_tile_m, init_tile_n = _l2_swizzle_tile(
+        bidx // cluster_m,
+        bidy // cluster_n,
+        gridx // cluster_m,
+        gridy // cluster_n,
+        swizzle_w,
+    )
+    init_tile_l = bidz
+
+    a_pattern = 0
+    for n_idx in cutlass.range_constexpr(cluster_n):
+        a_pattern = a_pattern | (1 << (n_idx * cluster_m))
+    b_pattern = 0
+    for pm_idx in cutlass.range_constexpr(cluster_m // 2):
+        b_pattern = b_pattern | (1 << (pm_idx * 2))
+
+    if cutlass.const_expr(multicast_a):
+        tma_mcast_mask_a = cutlass.Int16(a_pattern << m_rank)
+    else:
+        tma_mcast_mask_a = cutlass.Int16(1 << cta_rank_in_cluster)
+    if cutlass.const_expr(multicast_b):
+        tma_mcast_mask_b = cutlass.Int16((b_pattern << pair_member) << (n_rank * cluster_m))
+    else:
+        tma_mcast_mask_b = cutlass.Int16(1 << cta_rank_in_cluster)
+
+    _smem_sys_reserved = cutlass.Array(cutlass.Int8, 1024, space=cutlass.AddressSpace.smem, alignment=1)
+
+    ab_full_mbar_ptr = cutlass.Array(cutlass.Int64, ab_stages, space=cutlass.AddressSpace.smem)
+    ab_empty_mbar_ptr = cutlass.Array(cutlass.Int64, ab_stages, space=cutlass.AddressSpace.smem)
+    acc_empty_mbar_ptr = cutlass.Array(cutlass.Int64, acc_stages, space=cutlass.AddressSpace.smem)
+    acc_full_mbar_ptr = cutlass.Array(cutlass.Int64, acc_stages, space=cutlass.AddressSpace.smem)
+    tmem_dealloc_mbar_ptr = cutlass.Array(cutlass.Int64, 1, space=cutlass.AddressSpace.smem)
+    tmem_ptr_i32 = cutlass.Array(cutlass.Int32, 1, space=cutlass.AddressSpace.smem)
+
+    _clc_response_raw = cutlass.Array(cutlass.Int128, CLC_SCHED_STAGES, space=cutlass.AddressSpace.smem, alignment=16)
+    clc_response_ptr_base = cute.make_ptr(
+        cutlass.Int128,
+        _clc_response_raw.data_ptr(),
+        mem_space=cute.AddressSpace.smem,
+    )
+    clc_full_mbar_ptr = cutlass.Array(cutlass.Int64, CLC_SCHED_STAGES, space=cutlass.AddressSpace.smem, alignment=8)
+    clc_empty_mbar_ptr = cutlass.Array(cutlass.Int64, CLC_SCHED_STAGES, space=cutlass.AddressSpace.smem, alignment=8)
+    clc_full_mbar_cute_base = cute.make_ptr(
+        cutlass.Int64,
+        clc_full_mbar_ptr.data_ptr(),
+        mem_space=cute.AddressSpace.smem,
+    )
+
+    sA_elems = sA_packed_elems
+    sB_elems = sB_packed_elems
+    smem_a_list = [
+        cutlass.Array(
+            ab_dtype,
+            sA_elems * ab_stages,
+            space=cutlass.AddressSpace.smem,
+            alignment=1024,
+        )
+        for _ in range(num_a_operands)
+    ]
+    smem_b_list = [
+        cutlass.Array(
+            ab_dtype,
+            sB_elems * ab_stages,
+            space=cutlass.AddressSpace.smem,
+            alignment=1024,
+        )
+        for _ in range(num_b_operands)
+    ]
+    smem_sfa_list = [
+        cutlass.Array(
+            cutlass.Uint8,
+            sfa_smem_bytes * ab_stages,
+            space=cutlass.AddressSpace.smem,
+            alignment=1024,
+        )
+        for _ in range(num_a_operands)
+    ]
+    smem_sfb_list = [
+        cutlass.Array(
+            cutlass.Uint8,
+            sfb_smem_bytes * ab_stages,
+            space=cutlass.AddressSpace.smem,
+            alignment=1024,
+        )
+        for _ in range(num_b_operands)
+    ]
+
+    # @@TMA_STORE_ONLY:BEGIN@@
+    epi_subtile_elems = epi_tile_mn[0] * epi_tile_mn[1]
+    smem_d_ptr = cutlass.Array(
+        cd_dtype,
+        epi_subtile_elems * EPI_SMEM_STAGES,
+        space=cutlass.AddressSpace.smem,
+        alignment=1024,
+    )
+    # @@TMA_STORE_ONLY:END@@
+
+    acc_empty_count = num_epilogue_warps * 2
+    cta_group = 2
+    ab_empty_count = (cluster_m // cta_group) + cluster_n - 1
+    num_consumer_warps_per_cta = 7
+    clc_empty_count = num_consumer_warps_per_cta * cluster_size
+    if warp_idx == 0:
+        if cutlass.const_expr(use_acc_overlap):
+            if elect_one:
+                nvvm.mbarrier_init(tmem_dealloc_mbar_ptr, num_epilogue_warps)
+        else:
+            if elect_one:
+                nvvm.mbarrier_init(tmem_dealloc_mbar_ptr, 32)
+        for i in range(ab_stages):
+            if elect_one:
+                nvvm.mbarrier_init(ab_full_mbar_ptr.subview(i), 1)
+            if elect_one:
+                nvvm.mbarrier_init(ab_empty_mbar_ptr.subview(i), ab_empty_count)
+        for i in range(acc_stages):
+            if elect_one:
+                nvvm.mbarrier_init(acc_full_mbar_ptr.subview(i), 1)
+            if elect_one:
+                nvvm.mbarrier_init(acc_empty_mbar_ptr.subview(i), acc_empty_count)
+        for i in range(CLC_SCHED_STAGES):
+            if elect_one:
+                nvvm.mbarrier_init(clc_full_mbar_ptr.subview(i), 1)
+            if elect_one:
+                nvvm.mbarrier_init(clc_empty_mbar_ptr.subview(i), clc_empty_count)
+    nvvm.fence_mbarrier_init()
+    nvvm.barrier_cluster_arrive_relaxed()
+
+    sA_bytes = sA_elems * (ab_dtype.width // 8)
+    sB_bytes = sB_elems * (ab_dtype.width // 8)
+    num_tma_copy_bytes = (num_a_operands * (sA_bytes + sfa_smem_bytes) + num_b_operands * (sB_bytes + sfb_smem_bytes)) * 2
+
+    pair_n_size = cgrp_tile_mnk[1] // cluster_n
+    # Per-CTA output rows one MMA-M block covers. The pair splits M, so this is
+    # the per-CTA mma_inst_m — half the instruction's hardware M.
+    epi_rows_per_mma_m = cta_tile_mnk[0] // num_mma_m
+    if cutlass.const_expr(epi_rows_per_mma_m == 64):
+        # cluster-MMA m=128: the pair also splits N, so each CTA drains N/2.
+        cols_per_acc_stage = pair_n_size // 2
+    else:
+        cols_per_acc_stage = pair_n_size
+    tmem_alloc_bar_count = (num_epilogue_warps + 1) * 32
+
+    nvvm.barrier_cluster_wait()
+    nvvm.barrier_cta_sync(0)
+
+    # @@INJECT_TAP_PTRS@@
+
+    VEC_BYTES = vec_bytes_epi
+    vsize = (VEC_BYTES * 8) // cd_dtype.width
+
+    M = m
+    N = n
+    num_tile_m = cute.ceil_div(M, cgrp_tile_mnk[0])
+    num_tile_n = cute.ceil_div(N, cgrp_tile_mnk[1])
+    num_k_tiles = cute.ceil_div(k, cgrp_tile_mnk[2])
+
+    if warp_idx == scheduler_warp_id:
+        nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
+        sched_iter = cutlass.Int32(0)
+        clc_empty_phase = cutlass.Int32(1)
+        clc_full_phase = cutlass.Int32(0)
+        is_valid_sched = cutlass.Int32(1)
+        while is_valid_sched != 0:
+            stage = sched_iter % CLC_SCHED_STAGES
+            if stage == 0 and sched_iter != 0:
+                clc_empty_phase = clc_empty_phase ^ 1
+                clc_full_phase = clc_full_phase ^ 1
+
+            if is_cluster_leader_cta:
+                while not nvvm.mbarrier_try_wait_parity(clc_empty_mbar_ptr.subview(stage), clc_empty_phase, time_limit=10_000_000):
+                    pass
+
+            if elect_one:
+                nvvm.mbarrier_arrive_expect_tx(clc_full_mbar_ptr.subview(stage), 16)
+
+            if is_cluster_leader_cta:
+                if elect_one:
+                    cute_clc.issue_clc_query(
+                        clc_full_mbar_cute_base + stage,
+                        clc_response_ptr_base + stage,
+                        multicast=True,
+                    )
+
+            while not nvvm.mbarrier_try_wait_parity(clc_full_mbar_ptr.subview(stage), clc_full_phase, time_limit=10_000_000):
+                pass
+
+            _m_idx, _n_idx, _l_idx, vld = cute_clc.clc_response(clc_response_ptr_base + stage)
+            cute.arch.fence_proxy("async.shared", space="cta")
+            is_valid_sched = vld
+
+            nvvm.bar_warp_sync(0xFFFFFFFF)
+            if elect_one:
+                empty_remote = nvvm.mapa(clc_empty_mbar_ptr.subview(stage), 0)
+                nvvm.mbarrier_arrive(empty_remote, scope=nvvm.MemScope.CLUSTER, relaxed=True)
+
+            sched_iter += 1
+
+        if cutlass.const_expr(cluster_shape_mnk[0] * cluster_shape_mnk[1] > 1):
+            if is_cluster_leader_cta:
+                for _ in range(CLC_SCHED_STAGES):
+                    stage = sched_iter % CLC_SCHED_STAGES
+                    if stage == 0 and sched_iter != 0:
+                        clc_empty_phase = clc_empty_phase ^ 1
+                    while not nvvm.mbarrier_try_wait_parity(
+                        clc_empty_mbar_ptr.subview(stage),
+                        clc_empty_phase,
+                        time_limit=10_000_000,
+                    ):
+                        pass
+                    sched_iter += 1
+
+    if warp_idx == tma_warp_id:
+        nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
+        if cutlass.const_expr(USE_PDL):
+            if elect_one:
+                nvvm.griddepcontrol("wait")
+        ab_empty_phase_bit = cutlass.Int32(1)
+        ab_iter = cutlass.Int32(0)
+        tile_m = init_tile_m
+        tile_n = init_tile_n
+        tile_l = init_tile_l
+        tile_iter = cutlass.Int32(0)
+        is_valid = cutlass.Int32(1)
+        clc_full_phase_tma = cutlass.Int32(0)
+        while is_valid != 0:
+            logical_cta_tile_n = cgrp_tile_mnk[1] // cluster_n
+            coord_m_per_cta = tile_m * cgrp_tile_mnk[0] + m_rank * cta_tile_mnk[0]
+            coord_n_per_cta = tile_n * cgrp_tile_mnk[1] + n_rank * logical_cta_tile_n + pair_member * cta_tile_mnk[1]
+            if cutlass.const_expr(matmul_a_batch == 1):
+                tile_l_a = cutlass.Int32(0)
+            else:
+                tile_l_a = tile_l
+            if cutlass.const_expr(matmul_b_batch == 1):
+                tile_l_b = cutlass.Int32(0)
+            else:
+                tile_l_b = tile_l
+
+            for k_tile_idx in range(num_k_tiles):
+                stage = ab_iter % ab_stages
+                if stage == 0 and ab_iter != 0:
+                    ab_empty_phase_bit = ab_empty_phase_bit ^ 1
+
+                while not nvvm.mbarrier_try_wait_parity(ab_empty_mbar_ptr.subview(stage), ab_empty_phase_bit, time_limit=10_000_000):
+                    pass
+
+                coord_k = k_tile_idx * cgrp_tile_mnk[2]
+                coord_sf_k = k_tile_idx * sf_tma_box_k
+                coord_n_pair = tile_n * cgrp_tile_mnk[1] + n_rank * logical_cta_tile_n
+                sfb_n_block = coord_n_pair // 128
+
+                if is_pair_leader:
+                    if elect_one:
+                        nvvm.mbarrier_arrive_expect_tx(ab_full_mbar_ptr.subview(stage), num_tma_copy_bytes)
+
+                for _ai in cutlass.range_constexpr(num_a_operands):
+                    sA_stage = smem_a_list[_ai].subview(sA_elems * stage)
+                    tma_a_desc = tma_a_descs[_ai]
+                    sSFA_stage = smem_sfa_list[_ai].subview(sfa_smem_bytes * stage)
+                    tma_sfa_desc = tma_sfa_descs[_ai]
+                    sfa_m_block = coord_m_per_cta // 128
+                    if cutlass.const_expr(multicast_a):
+                        if n_rank == 0:
+                            if cutlass.const_expr(a_is_m_major):
+                                for m_group in cutlass.range_constexpr(cta_tile_mnk[0] // a_tma_group_elems):
+                                    if elect_one:
+                                        nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                            sA_stage.subview(m_group * a_tma_group_elems * cgrp_tile_mnk[2]),
+                                            tma_a_desc.get_ptr(),
+                                            (
+                                                coord_m_per_cta + m_group * a_tma_group_elems,
+                                                coord_k,
+                                                tile_l_a,
+                                            ),
+                                            ab_full_mbar_ptr.subview(stage),
+                                            [],
+                                            multicast_mask=tma_mcast_mask_a,
+                                            group=nvvm.CTAGroup.CTA_2,
+                                        )
+                            else:
+                                if elect_one:
+                                    nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                        sA_stage,
+                                        tma_a_desc.get_ptr(),
+                                        (coord_k, coord_m_per_cta, tile_l_a),
+                                        ab_full_mbar_ptr.subview(stage),
+                                        [],
+                                        multicast_mask=tma_mcast_mask_a,
+                                        group=nvvm.CTAGroup.CTA_2,
+                                    )
+                    else:
+                        if cutlass.const_expr(a_is_m_major):
+                            for m_group in cutlass.range_constexpr(cta_tile_mnk[0] // a_tma_group_elems):
+                                if elect_one:
+                                    nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                        sA_stage.subview(m_group * a_tma_group_elems * cgrp_tile_mnk[2]),
+                                        tma_a_desc.get_ptr(),
+                                        (
+                                            coord_m_per_cta + m_group * a_tma_group_elems,
+                                            coord_k,
+                                            tile_l_a,
+                                        ),
+                                        ab_full_mbar_ptr.subview(stage),
+                                        [],
+                                        multicast_mask=tma_mcast_mask_a,
+                                        group=nvvm.CTAGroup.CTA_2,
+                                    )
+                        else:
+                            if elect_one:
+                                nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                    sA_stage,
+                                    tma_a_desc.get_ptr(),
+                                    (coord_k, coord_m_per_cta, tile_l_a),
+                                    ab_full_mbar_ptr.subview(stage),
+                                    [],
+                                    multicast_mask=tma_mcast_mask_a,
+                                    group=nvvm.CTAGroup.CTA_2,
+                                )
+                    if cutlass.const_expr(multicast_a):
+                        if n_rank == 0:
+                            if elect_one:
+                                nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                    sSFA_stage,
+                                    tma_sfa_desc.get_ptr(),
+                                    (0, coord_sf_k, sfa_m_block, tile_l_a),
+                                    ab_full_mbar_ptr.subview(stage),
+                                    [],
+                                    multicast_mask=tma_mcast_mask_a,
+                                    group=nvvm.CTAGroup.CTA_2,
+                                )
+                    else:
+                        if elect_one:
+                            nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                sSFA_stage,
+                                tma_sfa_desc.get_ptr(),
+                                (0, coord_sf_k, sfa_m_block, tile_l_a),
+                                ab_full_mbar_ptr.subview(stage),
+                                [],
+                                multicast_mask=tma_mcast_mask_a,
+                                group=nvvm.CTAGroup.CTA_2,
+                            )
+
+                for _bj in cutlass.range_constexpr(num_b_operands):
+                    sB_stage = smem_b_list[_bj].subview(sB_elems * stage)
+                    tma_b_desc = tma_b_descs[_bj]
+                    sSFB_stage = smem_sfb_list[_bj].subview(sfb_smem_bytes * stage)
+                    tma_sfb_desc = tma_sfb_descs[_bj]
+                    if cutlass.const_expr(multicast_b):
+                        if pair_m_idx == 0:
+                            if cutlass.const_expr(b_is_n_major):
+                                for n_group in cutlass.range_constexpr(cta_tile_mnk[1] // b_tma_group_elems):
+                                    if elect_one:
+                                        nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                            sB_stage.subview(n_group * b_tma_group_elems * cgrp_tile_mnk[2]),
+                                            tma_b_desc.get_ptr(),
+                                            (
+                                                coord_n_per_cta + n_group * b_tma_group_elems,
+                                                coord_k,
+                                                tile_l_b,
+                                            ),
+                                            ab_full_mbar_ptr.subview(stage),
+                                            [],
+                                            multicast_mask=tma_mcast_mask_b,
+                                            group=nvvm.CTAGroup.CTA_2,
+                                        )
+                            else:
+                                if elect_one:
+                                    nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                        sB_stage,
+                                        tma_b_desc.get_ptr(),
+                                        (coord_k, coord_n_per_cta, tile_l_b),
+                                        ab_full_mbar_ptr.subview(stage),
+                                        [],
+                                        multicast_mask=tma_mcast_mask_b,
+                                        group=nvvm.CTAGroup.CTA_2,
+                                    )
+                    else:
+                        if cutlass.const_expr(b_is_n_major):
+                            for n_group in cutlass.range_constexpr(cta_tile_mnk[1] // b_tma_group_elems):
+                                if elect_one:
+                                    nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                        sB_stage.subview(n_group * b_tma_group_elems * cgrp_tile_mnk[2]),
+                                        tma_b_desc.get_ptr(),
+                                        (
+                                            coord_n_per_cta + n_group * b_tma_group_elems,
+                                            coord_k,
+                                            tile_l_b,
+                                        ),
+                                        ab_full_mbar_ptr.subview(stage),
+                                        [],
+                                        multicast_mask=tma_mcast_mask_b,
+                                        group=nvvm.CTAGroup.CTA_2,
+                                    )
+                        else:
+                            if elect_one:
+                                nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                    sB_stage,
+                                    tma_b_desc.get_ptr(),
+                                    (coord_k, coord_n_per_cta, tile_l_b),
+                                    ab_full_mbar_ptr.subview(stage),
+                                    [],
+                                    multicast_mask=tma_mcast_mask_b,
+                                    group=nvvm.CTAGroup.CTA_2,
+                                )
+                    if cutlass.const_expr(multicast_b):
+                        if pair_m_idx == 0:
+                            if elect_one:
+                                nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                    sSFB_stage,
+                                    tma_sfb_desc.get_ptr(),
+                                    (0, coord_sf_k, sfb_n_block, tile_l_b),
+                                    ab_full_mbar_ptr.subview(stage),
+                                    [],
+                                    multicast_mask=tma_mcast_mask_b,
+                                    group=nvvm.CTAGroup.CTA_2,
+                                )
+                    else:
+                        if elect_one:
+                            nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                sSFB_stage,
+                                tma_sfb_desc.get_ptr(),
+                                (0, coord_sf_k, sfb_n_block, tile_l_b),
+                                ab_full_mbar_ptr.subview(stage),
+                                [],
+                                multicast_mask=tma_mcast_mask_b,
+                                group=nvvm.CTAGroup.CTA_2,
+                            )
+
+                ab_iter += 1
+
+            consumer_stage = tile_iter % CLC_SCHED_STAGES
+            if consumer_stage == 0 and tile_iter != 0:
+                clc_full_phase_tma = clc_full_phase_tma ^ 1
+            while not nvvm.mbarrier_try_wait_parity(
+                clc_full_mbar_ptr.subview(consumer_stage),
+                clc_full_phase_tma,
+                time_limit=10_000_000,
+            ):
+                pass
+            m_idx, n_idx, l_idx, vld = cute_clc.clc_response(clc_response_ptr_base + consumer_stage)
+            cute.arch.fence_proxy("async.shared", space="cta")
+            is_valid = vld
+            tile_m, tile_n = _l2_swizzle_tile(
+                m_idx // cluster_m,
+                n_idx // cluster_n,
+                gridx // cluster_m,
+                gridy // cluster_n,
+                swizzle_w,
+            )
+            tile_l = l_idx
+            nvvm.bar_warp_sync(0xFFFFFFFF)
+            if elect_one:
+                empty_remote = nvvm.mapa(clc_empty_mbar_ptr.subview(consumer_stage), 0)
+                nvvm.mbarrier_arrive(empty_remote, scope=nvvm.MemScope.CLUSTER, relaxed=True)
+            tile_iter += 1
+
+        tail_stage = ab_iter % ab_stages
+        tail_phase = ab_empty_phase_bit
+        if tail_stage == 0 and ab_iter != 0:
+            tail_phase = tail_phase ^ 1
+        for _ in range(ab_stages - 1):
+            tail_stage = tail_stage + 1
+            if tail_stage == ab_stages:
+                tail_stage = cutlass.Int32(0)
+                tail_phase = tail_phase ^ 1
+        if elect_one:
+            while not nvvm.mbarrier_try_wait_parity(ab_empty_mbar_ptr.subview(tail_stage), tail_phase, time_limit=10_000_000):
+                pass
+
+    pair_mask = cutlass.Int16(3) << pair_leader_rank
+    a_arrive_pattern = 0
+    for n_idx in cutlass.range_constexpr(cluster_n):
+        a_arrive_pattern = a_arrive_pattern | (1 << (n_idx * cluster_m))
+    b_arrive_pattern = 0
+    for m_idx in cutlass.range_constexpr(cluster_m):
+        b_arrive_pattern = b_arrive_pattern | (1 << m_idx)
+    a_part = a_arrive_pattern << m_rank
+    a_part = a_part | (a_part << 1)
+    b_part = b_arrive_pattern << (n_rank * cluster_m)
+    ab_empty_arrive_mask = cutlass.Int16(a_part | b_part)
+    if warp_idx == mma_warp_id:
+        nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
+        nvvm.tcgen05_alloc(
+            tmem_ptr_i32,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_2,
+        )
+        nvvm.bar_warp_sync(0xFFFFFFFF)
+        nvvm.barrier_cta_arrive(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
+        tmem_raw_addr = tmem_ptr_i32.load()
+        base_col_id_root = tmem_raw_addr & 0xFFFF
+        base_row_id = tmem_raw_addr >> 16
+        peer_cta_rank = cta_rank_in_cluster ^ 1
+        if is_pair_leader:
+            ab_full_phase_bit = cutlass.Int32(0)
+            ab_iter = cutlass.Int32(0)
+            acc_empty_phase_bit = cutlass.Int32(1)
+            tile_iter = cutlass.Int32(0)
+            is_valid = cutlass.Int32(1)
+            clc_full_phase_mma = cutlass.Int32(0)
+            acc_stage = cutlass.Int32(0)
+
+            # fp4 packs its K-mode into the OMMA descriptor's 2-bit split field;
+            # fp8 keeps the MX descriptor's 1-bit one. Both are built once,
+            # outside the loops — the fields depend only on j (the scale id
+            # within a word).
+            if cutlass.const_expr(idesc_is_omma):
+                idesc_by_j = [
+                    cutlass.experimental.primitives.Tcgen05MxOmmaInstrDesc.build(
+                        a_dtype=idesc_a_dtype,
+                        b_dtype=idesc_b_dtype,
+                        scale_format=sf_scale_format,
+                        n_dim=mma_n_dim,
+                        m_dim=mma_m_dim,
+                        a_major=mma_a_major,
+                        b_major=mma_b_major,
+                        a_sf_id=j * sf_scales_per_inst,
+                        b_sf_id=j * sf_scales_per_inst,
+                        k_dim=mma_k_dim_mode,
+                    )
+                    for j in range(sf_insts_per_atom)
+                ]
+            else:
+                idesc_by_j = [
+                    cutlass.experimental.primitives.Tcgen05MxInstrDesc.build(
+                        a_dtype=idesc_a_dtype,
+                        b_dtype=idesc_b_dtype,
+                        scale_format=sf_scale_format,
+                        n_dim=mma_n_dim,
+                        m_dim=mma_m_dim,
+                        a_major=mma_a_major,
+                        b_major=mma_b_major,
+                        a_sf_id=j * sf_scales_per_inst,
+                        b_sf_id=j * sf_scales_per_inst,
+                        k_dim=mma_k_dim_mode,
+                    )
+                    for j in range(sf_insts_per_atom)
+                ]
+
+            sfa_tmem_bases = [(base_row_id << 16) | (base_col_id_root + sfa_col_bases[i]) for i in range(num_a_operands)]
+            sfb_tmem_bases = [(base_row_id << 16) | (base_col_id_root + sfb_col_bases[j]) for j in range(num_b_operands)]
+            s2t_shape, s2t_multicast = nvvm.S2TCopyMode.S2T_32x128b_WARPX4
+            sfa_scale_ptrs = [nvvm.make_tmem_ptr(b, cutlass.Float32) for b in sfa_tmem_bases]
+            sfb_scale_ptrs = [nvvm.make_tmem_ptr(b, cutlass.Float32) for b in sfb_tmem_bases]
+            # utccp destination per (MN-block, atom within the scale word). SFB
+            # is atom-MAJOR across the N-blocks because ONE instruction walks
+            # all of them; SFA is block-major because one instruction covers
+            # exactly one 128-row block, so that word has to be contiguous.
+            # Both collapse to the same addresses at a single block, and to
+            # sm100's layout at word_atoms == 1.
+            sfa_dst_ptrs = [
+                [
+                    [nvvm.make_tmem_ptr(sfa_tmem_bases[i] + m * registers_per_block + a * registers_per_atom, cutlass.Float32) for a in range(word_atoms)]
+                    for m in range(num_blocks_m)
+                ]
+                for i in range(num_a_operands)
+            ]
+            sfb_dst_ptrs = [
+                [
+                    [nvvm.make_tmem_ptr(sfb_tmem_bases[j] + (a * num_blocks_n + m) * registers_per_atom, cutlass.Float32) for a in range(word_atoms)]
+                    for m in range(num_blocks_n)
+                ]
+                for j in range(num_b_operands)
+            ]
+            while is_valid != 0:
+                acc_stage = tile_iter % acc_stages
+                if acc_stage == 0 and tile_iter != 0:
+                    acc_empty_phase_bit = acc_empty_phase_bit ^ 1
+
+                while not nvvm.mbarrier_try_wait_parity(
+                    acc_empty_mbar_ptr.subview(acc_stage),
+                    acc_empty_phase_bit,
+                    time_limit=10_000_000,
+                ):
+                    pass
+
+                if cutlass.const_expr(use_acc_overlap):
+                    acc_base_col = base_col_id_root + (tile_iter % 2) * acc_stage_stride
+                else:
+                    acc_base_col = base_col_id_root + acc_stage * acc_region_cols
+                # One accumulator per (gemm, M block); M block mi sits
+                # epi_cols_per_mma_m columns further into its GEMM's region and
+                # reads SF word block mi (SF words are one per 128 rows).
+                acc_tmem_ptrs = [
+                    [
+                        nvvm.make_tmem_ptr(
+                            (base_row_id << 16) | (acc_base_col + g * acc_gemm_stride + mi * epi_cols_per_mma_m),
+                            cutlass.Float32,
+                        )
+                        for mi in range(num_mma_m)
+                    ]
+                    for g in range(num_gemms)
+                ]
+
+                scale_d = cutlass.Boolean(False)
+                for k_tile_idx in range(num_k_tiles):
+                    stage = ab_iter % ab_stages
+                    if stage == 0 and ab_iter != 0:
+                        ab_full_phase_bit = ab_full_phase_bit ^ 1
+
+                    while not nvvm.mbarrier_try_wait_parity(
+                        ab_full_mbar_ptr.subview(stage),
+                        ab_full_phase_bit,
+                        time_limit=10_000_000,
+                    ):
+                        pass
+
+                    desc_a_bases = [
+                        cutlass.experimental.primitives.Tcgen05SmemDesc.build(
+                            start_address=smem_a_list[i].subview(sA_elems * stage),
+                            leading_byte_offset=a_smem_desc_leading_byte_offset,
+                            stride_byte_offset=a_smem_desc_stride_byte_offset,
+                            layout=ab_smem_swizzle,
+                        )
+                        for i in range(num_a_operands)
+                    ]
+                    desc_b_bases = [
+                        cutlass.experimental.primitives.Tcgen05SmemDesc.build(
+                            start_address=smem_b_list[j].subview(sB_elems * stage),
+                            leading_byte_offset=b_smem_desc_leading_byte_offset,
+                            stride_byte_offset=b_smem_desc_stride_byte_offset,
+                            layout=ab_smem_swizzle,
+                        )
+                        for j in range(num_b_operands)
+                    ]
+                    desc_sfa_bases = [
+                        cutlass.experimental.primitives.Tcgen05SmemDesc.build(
+                            start_address=smem_sfa_list[i].subview(sfa_smem_bytes * stage),
+                            leading_byte_offset=16,
+                            stride_byte_offset=128,
+                            layout=cutlass.experimental.primitives.Tcgen05SmemSwizzle.NONE,
+                        )
+                        for i in range(num_a_operands)
+                    ]
+                    desc_sfb_bases = [
+                        cutlass.experimental.primitives.Tcgen05SmemDesc.build(
+                            start_address=smem_sfb_list[j].subview(sfb_smem_bytes * stage),
+                            leading_byte_offset=16,
+                            stride_byte_offset=128,
+                            layout=cutlass.experimental.primitives.Tcgen05SmemSwizzle.NONE,
+                        )
+                        for j in range(num_b_operands)
+                    ]
+
+                    # One SF word per group of MMAs, refreshed right before they
+                    # read it. A word spans word_atoms consecutive K-atoms in SMEM.
+                    for atom_r in cutlass.range(num_sf_atoms, unroll_full=True):
+                        for _ai in cutlass.range_constexpr(num_a_operands):
+                            for _m in cutlass.range_constexpr(num_blocks_m):
+                                for _a in cutlass.range_constexpr(word_atoms):
+                                    if elect_one:
+                                        nvvm.tcgen05_cp(
+                                            s2t_shape,
+                                            sfa_dst_ptrs[_ai][_m][_a],
+                                            desc_sfa_bases[_ai] + (sf_atom_desc_stride * (atom_r * word_atoms + _a) + sf_block_desc_stride * _m),
+                                            group=nvvm.CTAGroup.CTA_2,
+                                            multicast=s2t_multicast,
+                                        )
+                        for _bj in cutlass.range_constexpr(num_b_operands):
+                            for _m in cutlass.range_constexpr(num_blocks_n):
+                                for _a in cutlass.range_constexpr(word_atoms):
+                                    if elect_one:
+                                        nvvm.tcgen05_cp(
+                                            s2t_shape,
+                                            sfb_dst_ptrs[_bj][_m][_a],
+                                            desc_sfb_bases[_bj] + (sf_atom_desc_stride * (atom_r * word_atoms + _a) + sf_block_desc_stride * _m),
+                                            group=nvvm.CTAGroup.CTA_2,
+                                            multicast=s2t_multicast,
+                                        )
+                        for j in cutlass.range_constexpr(sf_insts_per_atom):
+                            k_block_idx = atom_r * sf_insts_per_atom + j
+                            idesc_k = idesc_by_j[j]
+                            for g in cutlass.range_constexpr(num_gemms):
+                                _ai = gemm_a_idx[g]
+                                _bj = gemm_b_idx[g]
+                                desc_a_k = desc_a_bases[_ai].advance_start_address(a_smem_k_step_bytes * k_block_idx)
+                                desc_b = desc_b_bases[_bj].advance_start_address(b_smem_k_step_bytes * k_block_idx)
+                                for mi in cutlass.range_constexpr(num_mma_m):
+                                    # The M sub-block offset is a whole SMEM swizzle atom, so
+                                    # the descriptor's swizzle phase is preserved. B and its SF
+                                    # are shared; A's SF word block follows the M block.
+                                    desc_a = desc_a_k.advance_start_address(a_smem_m_step_bytes * mi)
+                                    if elect_one:
+                                        nvvm.tcgen05_mma_block_scale(
+                                            mma_block_scale_kind,
+                                            nvvm.CTAGroup.CTA_2,
+                                            acc_tmem_ptrs[g][mi],
+                                            desc_a,
+                                            desc_b,
+                                            idesc_k,
+                                            enable_input_d=scale_d,
+                                            scale_a=sfa_dst_ptrs[_ai][mi][0],
+                                            scale_b=sfb_scale_ptrs[_bj],
+                                            scale_vec_size=scale_vec_size,
+                                            b_collector_op=_b_collector_op(mi),
+                                        )
+                            # Every accumulator sees scale_d=False on exactly the first
+                            # k_block of the tile, so the flip stays outside mi.
+                            scale_d = cutlass.Boolean(True)
+
+                    if elect_one:
+                        nvvm.tcgen05_commit(
+                            ab_empty_mbar_ptr.subview(stage),
+                            multicast_mask=ab_empty_arrive_mask,
+                            group=nvvm.CTAGroup.CTA_2,
+                        )
+                    ab_iter += 1
+
+                if elect_one:
+                    nvvm.tcgen05_commit(
+                        acc_full_mbar_ptr.subview(acc_stage),
+                        multicast_mask=pair_mask,
+                        group=nvvm.CTAGroup.CTA_2,
+                    )
+
+                consumer_stage = tile_iter % CLC_SCHED_STAGES
+                if consumer_stage == 0 and tile_iter != 0:
+                    clc_full_phase_mma = clc_full_phase_mma ^ 1
+                while not nvvm.mbarrier_try_wait_parity(
+                    clc_full_mbar_ptr.subview(consumer_stage),
+                    clc_full_phase_mma,
+                    time_limit=10_000_000,
+                ):
+                    pass
+                _m_idx, _n_idx, _l_idx, vld = cute_clc.clc_response(clc_response_ptr_base + consumer_stage)
+                cute.arch.fence_proxy("async.shared", space="cta")
+                is_valid = vld
+                nvvm.bar_warp_sync(0xFFFFFFFF)
+                if elect_one:
+                    empty_remote = nvvm.mapa(clc_empty_mbar_ptr.subview(consumer_stage), 0)
+                    nvvm.mbarrier_arrive(empty_remote, scope=nvvm.MemScope.CLUSTER, relaxed=True)
+                tile_iter += 1
+
+            if cutlass.const_expr(USE_PDL):
+                if elect_one:
+                    nvvm.griddepcontrol("launch_dependents")
+
+            tail_stage = acc_stage
+            tail_phase = acc_empty_phase_bit
+            if elect_one:
+                for _ in range(acc_stages):
+                    tail_stage = tail_stage + 1
+                    if tail_stage == acc_stages:
+                        tail_stage = cutlass.Int32(0)
+                        tail_phase = tail_phase ^ 1
+                    while not nvvm.mbarrier_try_wait_parity(
+                        acc_empty_mbar_ptr.subview(tail_stage),
+                        tail_phase,
+                        time_limit=10_000_000,
+                    ):
+                        pass
+            nvvm.bar_warp_sync(0xFFFFFFFF)
+
+            nvvm.tcgen05_relinquish_alloc_permit(group=nvvm.CTAGroup.CTA_2)
+            peer_mbar = nvvm.mapa(tmem_dealloc_mbar_ptr, peer_cta_rank)
+            while not nvvm.mbarrier_try_wait_parity(tmem_dealloc_mbar_ptr, 0, time_limit=10_000_000):
+                pass
+            if cutlass.const_expr(not use_acc_overlap):
+                nvvm.mbarrier_arrive(peer_mbar, scope=nvvm.MemScope.CLUSTER, relaxed=True)
+            alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
+            nvvm.tcgen05_dealloc(
+                alloc_ptr,
+                cutlass.Int32(num_tmem_alloc_cols),
+                is_exclusive=tmem_alloc_exclusive,
+                group=nvvm.CTAGroup.CTA_2,
+            )
+        else:
+            tile_iter = cutlass.Int32(0)
+            is_valid = cutlass.Int32(1)
+            clc_full_phase_mma = cutlass.Int32(0)
+            while is_valid != 0:
+                consumer_stage = tile_iter % CLC_SCHED_STAGES
+                if consumer_stage == 0 and tile_iter != 0:
+                    clc_full_phase_mma = clc_full_phase_mma ^ 1
+                while not nvvm.mbarrier_try_wait_parity(
+                    clc_full_mbar_ptr.subview(consumer_stage),
+                    clc_full_phase_mma,
+                    time_limit=10_000_000,
+                ):
+                    pass
+                _m_idx, _n_idx, _l_idx, vld = cute_clc.clc_response(clc_response_ptr_base + consumer_stage)
+                cute.arch.fence_proxy("async.shared", space="cta")
+                is_valid = vld
+                nvvm.bar_warp_sync(0xFFFFFFFF)
+                if elect_one:
+                    empty_remote = nvvm.mapa(clc_empty_mbar_ptr.subview(consumer_stage), 0)
+                    nvvm.mbarrier_arrive(empty_remote, scope=nvvm.MemScope.CLUSTER, relaxed=True)
+                tile_iter += 1
+
+            if cutlass.const_expr(USE_PDL):
+                if elect_one:
+                    nvvm.griddepcontrol("launch_dependents")
+
+            nvvm.tcgen05_relinquish_alloc_permit(group=nvvm.CTAGroup.CTA_2)
+            peer_mbar = nvvm.mapa(tmem_dealloc_mbar_ptr, peer_cta_rank)
+            if cutlass.const_expr(not use_acc_overlap):
+                nvvm.mbarrier_arrive(peer_mbar, scope=nvvm.MemScope.CLUSTER, relaxed=True)
+            while not nvvm.mbarrier_try_wait_parity(tmem_dealloc_mbar_ptr, 0, time_limit=10_000_000):
+                pass
+            alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
+            nvvm.tcgen05_dealloc(
+                alloc_ptr,
+                cutlass.Int32(num_tmem_alloc_cols),
+                is_exclusive=tmem_alloc_exclusive,
+                group=nvvm.CTAGroup.CTA_2,
+            )
+
+    if warp_idx < num_epilogue_warps:
+        nvvm.setmaxregister(epi_reg_count, nvvm.SetMaxRegisterAction.INCREASE)
+        nvvm.barrier_cta_sync(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
+        tmem_raw_addr = tmem_ptr_i32.load()
+        base_col_id_root = tmem_raw_addr & 0xFFFF
+        base_row_id = tmem_raw_addr >> 16
+
+        if cutlass.const_expr(USE_PDL):
+            nvvm.griddepcontrol("wait")
+
+        tile_iter = cutlass.Int32(0)
+        acc_full_phase_bit = cutlass.Int32(0)
+        tile_m = init_tile_m
+        tile_n = init_tile_n
+        tile_l = init_tile_l
+        is_valid = cutlass.Int32(1)
+        clc_full_phase_epi = cutlass.Int32(0)
+
+        row_id_with_warp_offset = base_row_id + warp_idx * 32
+        if cutlass.const_expr(cols_per_acc_stage >= 32):
+            t2r_inst_repx = 32
+            subtile_cnt = cols_per_acc_stage // 32
+        else:
+            t2r_inst_repx = cols_per_acc_stage
+            subtile_cnt = 1
+        shape = nvvm.Tcgen05LdStShape.SHAPE_32X32B
+        lane = tidx % 32
+
+        # @@TMA_STORE_ONLY:BEGIN@@
+        epi_stage_idx = cutlass.Int32(EPI_SMEM_STAGES - 1)
+        # @@TMA_STORE_ONLY:END@@
+
+        while is_valid != 0:
+            coord_m_tile = tile_m * cgrp_tile_mnk[0] + m_rank * cta_tile_mnk[0]
+            coord_n_c = tile_n * cgrp_tile_mnk[1] + n_rank * pair_n_size
+            if cutlass.const_expr(epi_rows_per_mma_m == 64):
+                coord_n_c = coord_n_c + (warp_idx // 2) * cols_per_acc_stage
+
+            acc_stage = tile_iter % acc_stages
+            if acc_stage == 0 and tile_iter != 0:
+                acc_full_phase_bit = acc_full_phase_bit ^ 1
+
+            while not nvvm.mbarrier_try_wait_parity(acc_full_mbar_ptr.subview(acc_stage), acc_full_phase_bit, time_limit=10_000_000):
+                pass
+
+            if cutlass.const_expr(use_acc_overlap):
+                acc_buf_parity = tile_iter % 2
+                acc_base_col = base_col_id_root + acc_buf_parity * acc_stage_stride
+            else:
+                acc_buf_parity = cutlass.Int32(0)
+                acc_base_col = base_col_id_root + acc_stage * acc_region_cols
+            # The 2-CTA epilogue drains its own half of the instruction's M,
+            # epi_rows_per_mma_m rows at a time, so a CTA tile of num_mma_m blocks
+            # drains in num_mma_m passes over its own column region.
+            for mi in cutlass.range_constexpr(num_mma_m):
+                coord_m = coord_m_tile + mi * epi_rows_per_mma_m
+                mi_col_base = acc_base_col + mi * epi_cols_per_mma_m
+                tmem_col_addr_gemms = [(row_id_with_warp_offset << 16) | (mi_col_base + g * acc_gemm_stride) for g in range(num_gemms)]
+
+                if cutlass.const_expr(epi_rows_per_mma_m == 64):
+                    row = coord_m + (warp_idx % 2) * 32 + lane
+                    row_active = True
+                else:
+                    row = coord_m + tidx
+                    row_active = True
+
+                # @@INJECT_AUX_VIEWS@@
+
+                for subtile_idx in cutlass.range(subtile_cnt, unroll_full=True):
+                    if cutlass.const_expr(use_acc_overlap):
+                        _sub = subtile_idx + (1 - acc_buf_parity) * (subtile_cnt - 1 - 2 * subtile_idx)
+                        subtile_col_offset = _sub * t2r_inst_repx
+                    else:
+                        subtile_col_offset = subtile_idx * t2r_inst_repx
+
+                    if cutlass.const_expr(not (use_tma_store_epi and cd_out_is_m_major)):
+                        c_rmem_vecs = []
+                        for g in cutlass.range_constexpr(num_gemms):
+                            tmem = cutlass.inttoptr(
+                                tmem_col_addr_gemms[g] + subtile_col_offset,
+                                6,
+                                cutlass.Float32,
+                            )
+                            c_rmem_vecs.append(nvvm.tcgen05_ld(shape, tmem, num=t2r_inst_repx))
+                        c_rmem_vec = c_rmem_vecs[0]
+
+                    if use_acc_overlap and (not cd_out_is_m_major) and mi == num_mma_m - 1 and subtile_idx == acc_overlap_subtiles - 1:
+                        nvvm.tcgen05_wait(kind=nvvm.Tcgen05Wait.LOAD)
+                        nvvm.tcgen05_fence(nvvm.Tcgen05Fence.BEFORE_THREAD_SYNC)
+                        if elect_one:
+                            mbar_pair_ptr = nvvm.mapa(acc_empty_mbar_ptr.subview(acc_stage), pair_leader_rank)
+                            nvvm.mbarrier_arrive(mbar_pair_ptr, scope=nvvm.MemScope.CLUSTER, relaxed=True)
+
+                    col = coord_n_c + subtile_col_offset
+
+                    # @@TMA_STORE_ONLY:BEGIN@@
+                    epi_stage_idx = (epi_stage_idx + 1) % EPI_SMEM_STAGES
+                    smem_subtile_ptr = smem_d_ptr.subview(epi_stage_idx * epi_subtile_elems)
+                    smem_thr_ptr = smem_subtile_ptr.subview(tidx * t2r_inst_repx)
+
+                    if cutlass.const_expr(cd_out_is_m_major):
+                        ld_col = mi_col_base + subtile_col_offset
+                        for _h in cutlass.range(2, unroll_full=True):
+                            ld_row = base_row_id + warp_idx * 32 + _h * 16
+                            ld_addr = (ld_row << 16) | ld_col
+                            ld_tmem = cutlass.inttoptr(ld_addr, 6, cutlass.Float32)
+                            _lv = nvvm.tcgen05_ld(nvvm.Tcgen05LdStShape.SHAPE_16X256B, ld_tmem, num=4)
+                            vec_f32 = _lv
+                            col_j = col
+                            linear_idx = tile_l * out_stride_l_0 + row * out_stride_m_0 + col_j * out_stride_n_0
+
+                            # @@INJECT_EPILOGUE@@
+
+                            _i32 = vec_out.bitcast(cutlass.Int32)
+                            for _blk in cutlass.range_constexpr(2):
+                                _regs = [_i32[_blk * 4 + _j] for _j in range(4)]
+                                _n_full = (lane % 8) + 8 * (lane // 16) + 16 * _blk
+                                _m_base = warp_idx * 32 + _h * 16 + 8 * ((lane // 8) % 2)
+                                _stm_off = (
+                                    (_m_base // cd_mmajor_atom_m) * (cd_mmajor_atom_m * epi_tile_mn[1])
+                                    + (_m_base % cd_mmajor_atom_m)
+                                    + _n_full * cd_mmajor_atom_m
+                                )
+                                nvvm.stmatrix(
+                                    _apply_smem_swizzle(
+                                        smem_subtile_ptr.data_ptr() + _stm_off,
+                                        cutlass.Swizzle(3, 4, 3),
+                                    ),
+                                    _regs,
+                                    nvvm.MMALayout.COL,
+                                    shape=nvvm.StoreShape.M8N8,
+                                )
+                    else:
+                        vec_f32 = c_rmem_vec
+                        col_j = col
+                        linear_idx = tile_l * out_stride_l_0 + row * out_stride_m_0 + col_j * out_stride_n_0
+
+                        # @@INJECT_EPILOGUE@@
+
+                        smem_thr_ptr.data_ptr().store_swizzled(vec_out, alignment=64, swizzle=cutlass.Swizzle(2, 4, 3))
+
+                    cute.arch.fence_view_async_shared()
+                    nvvm.barrier_cta_sync(
+                        barrier_id=EPI_SYNC_BAR_ID,
+                        thread_count=num_epilogue_warps * 32,
+                    )
+
+                    if warp_idx == 0:
+                        if cutlass.const_expr(cd_out_is_m_major):
+                            for _mb in cutlass.range_constexpr(epi_tile_mn[0] // cd_mmajor_atom_m):
+                                if elect_one:
+                                    nvvm.cp_async_bulk_tensor_global_shared_cta(
+                                        tma_c_desc.get_ptr(),
+                                        smem_subtile_ptr.subview(_mb * (cd_mmajor_atom_m * epi_tile_mn[1])),
+                                        (coord_m + _mb * cd_mmajor_atom_m, col, tile_l),
+                                    )
+                        else:
+                            if elect_one:
+                                nvvm.cp_async_bulk_tensor_global_shared_cta(
+                                    tma_c_desc.get_ptr(),
+                                    smem_subtile_ptr,
+                                    (col, coord_m, tile_l),
+                                )
+                        if elect_one:
+                            nvvm.cp_async_bulk_commit_group()
+                        nvvm.cp_async_bulk_wait_group(EPI_SMEM_STAGES - 1, read=True)
+
+                    nvvm.barrier_cta_sync(
+                        barrier_id=EPI_SYNC_BAR_ID,
+                        thread_count=num_epilogue_warps * 32,
+                    )
+                    # @@TMA_STORE_ONLY:END@@
+
+                    # @@STG_ONLY:BEGIN@@
+                    if row_active and row < M:
+                        for j in cutlass.range_constexpr(t2r_inst_repx // vsize):
+                            col_j = col + j * vsize
+                            if col_j + vsize <= N:
+                                vec_f32 = c_rmem_vec[j * vsize : (j + 1) * vsize]
+
+                                # @@INJECT_STG_VEC_BINDINGS@@
+
+                                # @@INJECT_EPILOGUE@@
+                    # @@STG_ONLY:END@@
+
+            if cutlass.const_expr((not use_acc_overlap) or cd_out_is_m_major):
+                nvvm.tcgen05_wait(kind=nvvm.Tcgen05Wait.LOAD)
+                nvvm.tcgen05_fence(nvvm.Tcgen05Fence.BEFORE_THREAD_SYNC)
+                if elect_one:
+                    mbar_pair_ptr = nvvm.mapa(acc_empty_mbar_ptr.subview(acc_stage), pair_leader_rank)
+                    nvvm.mbarrier_arrive(mbar_pair_ptr, scope=nvvm.MemScope.CLUSTER, relaxed=True)
+
+            consumer_stage = tile_iter % CLC_SCHED_STAGES
+            if consumer_stage == 0 and tile_iter != 0:
+                clc_full_phase_epi = clc_full_phase_epi ^ 1
+            while not nvvm.mbarrier_try_wait_parity(
+                clc_full_mbar_ptr.subview(consumer_stage),
+                clc_full_phase_epi,
+                time_limit=10_000_000,
+            ):
+                pass
+            m_idx, n_idx, l_idx, vld = cute_clc.clc_response(clc_response_ptr_base + consumer_stage)
+            cute.arch.fence_proxy("async.shared", space="cta")
+            is_valid = vld
+            tile_m, tile_n = _l2_swizzle_tile(
+                m_idx // cluster_m,
+                n_idx // cluster_n,
+                gridx // cluster_m,
+                gridy // cluster_n,
+                swizzle_w,
+            )
+            tile_l = l_idx
+            nvvm.bar_warp_sync(0xFFFFFFFF)
+            if elect_one:
+                empty_remote = nvvm.mapa(clc_empty_mbar_ptr.subview(consumer_stage), 0)
+                nvvm.mbarrier_arrive(empty_remote, scope=nvvm.MemScope.CLUSTER, relaxed=True)
+
+            tile_iter += 1
+
+        if cutlass.const_expr(use_acc_overlap):
+            nvvm.tcgen05_wait(kind=nvvm.Tcgen05Wait.LOAD)
+            nvvm.tcgen05_fence(nvvm.Tcgen05Fence.BEFORE_THREAD_SYNC)
+            if elect_one:
+                nvvm.mbarrier_arrive(tmem_dealloc_mbar_ptr)
+
+        # @@TMA_STORE_ONLY:BEGIN@@
+        if warp_idx == 0:
+            nvvm.cp_async_bulk_wait_group(0, read=True)
+        # @@TMA_STORE_ONLY:END@@
+
+    if warp_idx == unused_warp_id:
+        nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
+
+
+@cute.jit
+def _host(
+    problem_size: tuple,
+    # @@INJECT_HOST_AB_PARAMS@@
+    # @@INJECT_HOST_TAP_PARAMS@@
+    # @@INJECT_HOST_AUX_PARAMS@@
+    # @@TMA_STORE_ONLY:BEGIN@@
+    # @@INJECT_HOST_TMA_C_PARAMS@@
+    # @@TMA_STORE_ONLY:END@@
+    stream: _cuda.CUstream,
+) -> None:
+    # @@INJECT_HOST_AB_LISTS@@
+    m = problem_size[0]
+    n = problem_size[1]
+    k_sym = problem_size[2]
+    batch = problem_size[3]
+    a_stride_m = problem_size[4]
+    a_stride_k = problem_size[5]
+    a_stride_l = problem_size[6]
+    b_stride_n = problem_size[7]
+    b_stride_k = problem_size[8]
+    b_stride_l = problem_size[9]
+    # @@INJECT_HOST_REDUCTION_STRIDES@@
+
+    if cutlass.const_expr(matmul_a_batch == 1):
+        a_batch = 1
+    else:
+        a_batch = batch
+    if cutlass.const_expr(matmul_b_batch == 1):
+        b_batch = 1
+    else:
+        b_batch = batch
+
+    rest_k = ((k_sym // block_size) + 3) // 4
+    rest_m = (m + 127) // 128
+    rest_n = (n + 127) // 128
+    tma_a_desc_list = []
+    tma_sfa_desc_list = []
+    for _a_op, _sfa_op in zip(_a_operands, _sfa_operands):
+        if cutlass.const_expr(a_is_m_major):
+            tma_a_desc_list.append(
+                _tma.create_tensor_map_tiled(
+                    global_address=_a_op.iterator.toint(),
+                    dtype=ab_tma_desc_dtype,
+                    global_dims=[m, k_sym, a_batch],
+                    global_strides=[
+                        a_stride_k * ab_dtype.width // 128,
+                        a_stride_l * ab_dtype.width // 128,
+                    ],
+                    box_dims=[a_tma_group_elems, cta_tile_mnk[2], 1],
+                    swizzle=ab_tma_swizzle,
+                    tma_format=ab_tma_format,
+                )
+            )
+        else:
+            tma_a_desc_list.append(
+                _tma.create_tensor_map_tiled(
+                    global_address=_a_op.iterator.toint(),
+                    dtype=ab_tma_desc_dtype,
+                    global_dims=[k_sym, m, a_batch],
+                    global_strides=[
+                        a_stride_m * ab_dtype.width // 128,
+                        a_stride_l * ab_dtype.width // 128,
+                    ],
+                    box_dims=[cta_tile_mnk[2], cta_tile_mnk[0], 1],
+                    swizzle=ab_tma_swizzle,
+                    tma_format=ab_tma_format,
+                )
+            )
+        sfa_fp16_tensor = cute.make_tensor(
+            cute.recast_ptr(_sfa_op.iterator, dtype=cutlass.Float16),
+            cute.make_layout(
+                (256, rest_k, rest_m, batch),
+                stride=(
+                    1,
+                    256,
+                    cute.assume(256 * rest_k, 8),
+                    cute.assume(256 * rest_k * rest_m, 8),
+                ),
+            ),
+        )
+        tma_sfa_desc_list.append(
+            _tma.create_tensor_map_tiled_from_view(
+                sfa_fp16_tensor,
+                dtype=cutlass.Uint16,
+                box_dims=(256, sf_tma_box_k, sfa_tma_box_mn, 1),
+                stride_order=(0, 1, 2, 3),
+                swizzle=_tma.TensorMapSwizzle.none,
+            )
+        )
+    tma_b_desc_list = []
+    tma_sfb_desc_list = []
+    for _b_op, _sfb_op in zip(_b_operands, _sfb_operands):
+        if cutlass.const_expr(b_is_n_major):
+            tma_b_desc_list.append(
+                _tma.create_tensor_map_tiled(
+                    global_address=_b_op.iterator.toint(),
+                    dtype=ab_tma_desc_dtype,
+                    global_dims=[n, k_sym, b_batch],
+                    global_strides=[
+                        b_stride_k * ab_dtype.width // 128,
+                        b_stride_l * ab_dtype.width // 128,
+                    ],
+                    box_dims=[b_tma_group_elems, cta_tile_mnk[2], 1],
+                    swizzle=ab_tma_swizzle,
+                    tma_format=ab_tma_format,
+                )
+            )
+        else:
+            tma_b_desc_list.append(
+                _tma.create_tensor_map_tiled(
+                    global_address=_b_op.iterator.toint(),
+                    dtype=ab_tma_desc_dtype,
+                    global_dims=[k_sym, n, b_batch],
+                    global_strides=[
+                        b_stride_n * ab_dtype.width // 128,
+                        b_stride_l * ab_dtype.width // 128,
+                    ],
+                    box_dims=[cta_tile_mnk[2], cta_tile_mnk[1], 1],
+                    swizzle=ab_tma_swizzle,
+                    tma_format=ab_tma_format,
+                )
+            )
+        sfb_fp16_tensor = cute.make_tensor(
+            cute.recast_ptr(_sfb_op.iterator, dtype=cutlass.Float16),
+            cute.make_layout(
+                (256, rest_k, rest_n, batch),
+                stride=(
+                    1,
+                    256,
+                    cute.assume(256 * rest_k, 8),
+                    cute.assume(256 * rest_k * rest_n, 8),
+                ),
+            ),
+        )
+        tma_sfb_desc_list.append(
+            _tma.create_tensor_map_tiled_from_view(
+                sfb_fp16_tensor,
+                dtype=cutlass.Uint16,
+                box_dims=(256, sf_tma_box_k, sfb_tma_box_mn, 1),
+                stride_order=(0, 1, 2, 3),
+                swizzle=_tma.TensorMapSwizzle.none,
+            )
+        )
+
+    # @@TMA_STORE_ONLY:BEGIN@@
+    # @@INJECT_HOST_TMA_C_LISTS@@
+    c = _tma_c_outputs[0]
+    if cutlass.const_expr(cd_out_is_m_major):
+        tma_c_desc = _tma.create_tensor_map_tiled(
+            global_address=c.iterator.toint(),
+            dtype=cd_tma_dtype,
+            global_dims=[m, n, batch],
+            global_strides=[
+                out_stride_n_0 * cd_dtype.width // 128,
+                out_stride_l_0 * cd_dtype.width // 128,
+            ],
+            box_dims=[cd_mmajor_atom_m, epi_tile_mn[1], 1],
+            swizzle=(_tma.TensorMapSwizzle.s128b if cutlass.const_expr(use_tma_store_epi) else _tma.TensorMapSwizzle.none),
+        )
+    else:
+        tma_c_desc = _tma.create_tensor_map_tiled(
+            global_address=c.iterator.toint(),
+            dtype=cd_tma_dtype,
+            global_dims=[n, m, batch],
+            global_strides=[
+                out_stride_m_0 * cd_dtype.width // 128,
+                out_stride_l_0 * cd_dtype.width // 128,
+            ],
+            box_dims=[epi_tile_mn[1], epi_tile_mn[0], 1],
+            swizzle=(_tma.TensorMapSwizzle.s64b if cutlass.const_expr(use_tma_store_epi) else _tma.TensorMapSwizzle.none),
+        )
+    tma_c_desc_list = [tma_c_desc]
+    # @@TMA_STORE_ONLY:END@@
+
+    cluster_m = cluster_shape_mnk[0]
+    cluster_n = cluster_shape_mnk[1]
+    cgrp_tile_m = cgrp_tile_mnk[0]
+    cgrp_tile_n = cgrp_tile_mnk[1]
+    num_tile_m_host = (m + cgrp_tile_m - 1) // cgrp_tile_m
+    num_tile_n_host = (n + cgrp_tile_n - 1) // cgrp_tile_n
+    grid_x = num_tile_m_host * cluster_m
+    grid_y = num_tile_n_host * cluster_n
+    grid_shape = (grid_x, grid_y, batch)
+    _kernel(
+        problem_size[0],
+        problem_size[1],
+        problem_size[2],
+        # @@INJECT_HOST_KERNEL_DESC_PASS@@
+        # @@INJECT_HOST_TAP_PASS@@
+        # @@INJECT_HOST_REDUCTION_STRIDE_PASS@@
+        # @@INJECT_HOST_AUX_PASS@@
+        # @@TMA_STORE_ONLY:BEGIN@@
+        # @@INJECT_HOST_TMA_C_PASS@@
+        # @@TMA_STORE_ONLY:END@@
+    ).launch(
+        grid=grid_shape,
+        block=(threads_per_cta, 1, 1),
+        cluster=cluster_shape_mnk,
+        use_pdl=USE_PDL,
+        stream=stream,
+    )
+
+
+@lru_cache(maxsize=None)
+def compile() -> Callable:
+    out_vec_elems = vec_bytes_epi // (cd_dtype.width // 8)
+    ab_stride_elems = 128 // ab_dtype.width
+    sym_m = cute.sym_int64()
+    sym_n = cute.sym_int64(divisibility=out_vec_elems)
+    # K tails are supported: the K loop is ceil_div and the TMA descriptor's global K
+    # extent makes a partial box HW zero-filled. The only real K rule is the 16-byte
+    # TMA contiguous-extent one, already gated by _tma_alignment_reject.
+    sym_k = cute.sym_int64()
+    # Packed K extent: same reasoning as sym_k -- no CTA-tile multiple is required.
+    sym_kp = cute.sym_int64()
+    sym_l = cute.sym_int64()
+    if matmul_a_batch == 1:
+        sym_a_l = 1
+    else:
+        sym_a_l = sym_l
+    if matmul_b_batch == 1:
+        sym_b_l = 1
+    else:
+        sym_b_l = sym_l
+
+    def _make_fake_a():
+        return make_fake_compact_tensor(
+            a_fake_dtype,
+            (sym_m, sym_kp, sym_a_l),
+            stride_order=(0, 1, 2) if a_is_m_major else (1, 0, 2),
+            assumed_align=16,
+        )
+
+    def _make_fake_b():
+        return make_fake_compact_tensor(
+            b_fake_dtype,
+            (sym_n, sym_kp, sym_b_l),
+            stride_order=(0, 1, 2) if b_is_n_major else (1, 0, 2),
+            assumed_align=16,
+        )
+
+    # SF reaches the kernel as a base pointer only; the host rebuilds the
+    # F8_128x4 view from problem_size, so no SF mode carries a layout contract.
+    def _make_fake_sfa():
+        return cute.runtime.make_fake_tensor(
+            sf_cutlass_dtype,
+            (cute.sym_int64(), cute.sym_int64(), cute.sym_int64()),
+            stride=(cute.sym_int64(), cute.sym_int64(), cute.sym_int64()),
+            assumed_align=16,
+        )
+
+    def _make_fake_sfb():
+        return cute.runtime.make_fake_tensor(
+            sf_cutlass_dtype,
+            (cute.sym_int64(), cute.sym_int64(), cute.sym_int64()),
+            stride=(cute.sym_int64(), cute.sym_int64(), cute.sym_int64()),
+            assumed_align=16,
+        )
+
+    # @@TMA_STORE_ONLY:BEGIN@@
+    def _make_fake_c():
+        return make_fake_compact_tensor(
+            cd_dtype,
+            (sym_m, sym_n // cd_fake_n_div, sym_l),
+            stride_order=(0, 1, 2) if cd_out_is_m_major else (1, 0, 2),
+            assumed_align=16,
+        )
+
+    # @@INJECT_COMPILE_TMA_C_FAKES@@
+    # @@TMA_STORE_ONLY:END@@
+    # @@INJECT_COMPILE_AB_FAKES@@
+    # The operand's unit stride (m/n when MN-major, k when K-major) never reaches TMA, so it carries no 16B contract.
+    sym_a_stride_m = cute.sym_int64() if a_is_m_major else cute.sym_int64(divisibility=ab_stride_elems)
+    sym_a_stride_k = cute.sym_int64(divisibility=ab_stride_elems) if a_is_m_major else cute.sym_int64()
+    sym_a_stride_l = cute.sym_int64(divisibility=ab_stride_elems)
+    sym_b_stride_n = cute.sym_int64() if b_is_n_major else cute.sym_int64(divisibility=ab_stride_elems)
+    sym_b_stride_k = cute.sym_int64(divisibility=ab_stride_elems) if b_is_n_major else cute.sym_int64()
+    sym_b_stride_l = cute.sym_int64(divisibility=ab_stride_elems)
+    # @@INJECT_COMPILE_REDUCTION_STRIDE_DECLS@@
+    # @@INJECT_COMPILE_TAP_FAKES@@
+    problem_size = (
+        sym_m,
+        sym_n,
+        sym_k,
+        sym_l,
+        sym_a_stride_m,
+        sym_a_stride_k,
+        sym_a_stride_l,
+        sym_b_stride_n,
+        sym_b_stride_k,
+        sym_b_stride_l,
+        # @@INJECT_COMPILE_REDUCTION_STRIDE_SYMBOLS@@
+    )
+    # @@INJECT_COMPILE_AUX_FAKES@@
+    _fake_stream = make_fake_stream(use_tvm_ffi_env_stream=False)
+    return cute.compile(
+        _host,
+        problem_size,
+        # @@INJECT_COMPILE_AB_PASS@@
+        # @@INJECT_COMPILE_TAP_PASS@@
+        # @@INJECT_COMPILE_AUX_PASS@@
+        # @@TMA_STORE_ONLY:BEGIN@@
+        # @@INJECT_COMPILE_TMA_C_PASS@@
+        # @@TMA_STORE_ONLY:END@@
+        stream=_fake_stream,
+        options=frost_compile_options,
+    )
diff --git a/python/cudnn/gemm/frost/kernel_templates/sm107_moe_grouped_block_scale_matmul_fwd_1ctamma.py b/python/cudnn/gemm/frost/kernel_templates/sm107_moe_grouped_block_scale_matmul_fwd_1ctamma.py
new file mode 100644
index 000000000..8f9bcbfbf
--- /dev/null
+++ b/python/cudnn/gemm/frost/kernel_templates/sm107_moe_grouped_block_scale_matmul_fwd_1ctamma.py
@@ -0,0 +1,1366 @@
+# Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: MIT
+
+"""sm107 cta_group=1 MoE grouped block-scale matmul fwd.
+
+Per routed group g: ``out[fto[g]:fto[g+1]] = deq(token[range]) @ deq(weight[g % E]).T``,
+where token/weight are FP4/FP8 dequantized by per-block scale factors inside the
+MMA. A grouped persistent scheduler walks the routed groups and the TMA warp
+patches A's tensormap per group change. Supports nvfp4 / mxfp4 / mxfp8 (K-major);
+single-CTA MMA.
+
+The pipeline is the sm100 one; SM 10.7's block-scale MMA reads a **64-byte K**
+per instruction instead of 32, which shows up in exactly two places (both
+driven by injected constants, so the rest of the file stays in lockstep with
+``sm100_moe_grouped_block_scale_matmul_fwd_1ctamma.py``):
+
+  * half as many MMAs per K-tile, each consuming ``sf_scales_per_inst`` scales
+    (8 at K-block 16, 4 at 32 — it follows the BLOCK SIZE, not the scale
+    dtype). When that exceeds the 4 scales one 128x4 utccp atom holds, a scale
+    *word* spans ``word_atoms`` atoms, and the two SF regions then lay them
+    out DIFFERENTLY: SFB atom-major across its N-blocks, SFA block-major.
+    At K-block 32 ``word_atoms == 1`` (identical to sm100).
+  * fp4 rides the OMMA instruction descriptor (``Tcgen05MxOmmaInstrDesc``,
+    K-mode 2 = 128 fp4 elements); mxfp8 stays on ``Tcgen05MxInstrDesc``
+    (K-mode 1 = 64 fp8 elements). Both take the real operand dtype.
+
+The K-tile itself is unchanged (128 bytes), so the grouped scheduler, the
+per-group A-tensormap patch and the per-group-128-padded SF blob layout are
+byte-for-byte the sm100 ones.
+
+Warp layout (8 warps × 32 = 256 threads/CTA): epilogue 0-3 (warp 0 allocates
+TMEM), MMA 4, TMA 5, grouped scheduler 6, unused donor 7.
+"""
+
+from __future__ import annotations
+
+from functools import lru_cache
+from typing import Callable
+
+import cutlass.experimental.primitives as nvvm
+from cudnn.gemm.frost.kernel_templates._tile_helpers import (
+    copy_tensormap_to_workspace as _copy_tensormap_to_workspace,
+    fence_tensormap_acquire as _fence_tensormap_acquire,
+    fence_tensormap_release as _fence_tensormap_release,
+    moe_swizzle_tile as _moe_swizzle_tile,
+    replace_tensormap_global_address as _replace_tensormap_global_address,
+    replace_tensormap_global_dim_1 as _replace_tensormap_global_dim_1,
+    TENSOR_MAP_QWORDS,
+)
+import cutlass.experimental.cuda.tensor_map as _tma
+import cutlass
+import cutlass.cute as cute
+from cutlass.cute.runtime import make_fake_compact_tensor
+from cutlass.cute.runtime import make_fake_stream
+from cuda.bindings import driver as _cuda
+
+# A TMA tensormap is 128 bytes = 16 int64 qwords. The per-group A descriptor
+# replacement keeps a per-CTA SMEM copy, patches base/M-dim there, then publishes
+# it to the per-CTA GMEM workspace the TMA reads.
+# @@INJECT_TILE_CONSTANTS@@
+
+
+# Per-CTA scheduler ring (replaces CLC): 2 stages, 8 int32 slot words.
+SCHED_STAGES = 2
+SCHED_SLOT_WORDS = 8
+
+USE_PDL = True
+EPI_SMEM_STAGES = 2
+EPI_SYNC_BAR_ID = 1
+TMEM_ALLOC_BARRIER_ID = 2
+
+
+@cute.jit
+def _moe_auto_swizzle_w(group_rows, n, k, nt_n):
+    """N-super-block width for one routed group, resolved per group.
+
+    Same rule as the dense path, but the "M side" is THIS group's token slice, not the
+    whole token tensor: block along the shorter of (group tokens, expert weight), capped
+    by what L2 can hold onto. A group spanning one m-tile makes both orders identical.
+    """
+    if cutlass.const_expr(tile_swizzle_n > 0):
+        return tile_swizzle_n
+    budget = cutlass.Int64(swizzle_l2_budget_bytes)
+    row_bytes = (cutlass.Int64(ab_dtype.width) * k) // 8
+    cap = cutlass.max(budget // (row_bytes * cgrp_tile_mnk[1]), cutlass.Int64(1))
+    w = cutlass.min(cutlass.Int64(nt_n), cap)
+    rows = cutlass.Int64(group_rows)
+    if cutlass.min(rows, n) * row_bytes <= budget and rows <= n:
+        w = cutlass.Int64(1)
+    return cutlass.Int32(w)
+
+
+def _b_collector_op(mi):
+    """B is identical across the M sub-blocks (only A's address advances), so the
+    first MMA fills the B collector and the rest read it back instead of
+    re-fetching the same operand from SMEM."""
+    if cutlass.const_expr(not b_collector_ok or num_mma_m == 1):
+        return None
+    if cutlass.const_expr(mi == 0):
+        return nvvm.Tcgen05MMACollectorOp.FILL
+    if cutlass.const_expr(mi == num_mma_m - 1):
+        return nvvm.Tcgen05MMACollectorOp.LASTUSE
+    return nvvm.Tcgen05MMACollectorOp.USE
+
+
+@cute.kernel
+def _kernel(
+    m: cutlass.Int64,
+    n: cutlass.Int64,
+    k: cutlass.Int64,
+    num_experts: cutlass.Int32,
+    num_groups: cutlass.Int32,
+    first_token_offset: cute.Tensor,
+    a_tma_workspace: cute.Tensor,
+    # @@INJECT_KERNEL_AB_DESC_PARAMS@@
+    # @@INJECT_MOE_KERNEL_MA_PARAMS@@
+    # @@INJECT_KERNEL_TAP_PARAMS@@
+    # @@INJECT_KERNEL_REDUCTION_STRIDE_PARAMS@@
+    # @@INJECT_KERNEL_AUX_PARAMS@@
+) -> None:
+    # @@INJECT_AB_DESC_LISTS@@
+    # @@INJECT_MOE_MA_LIST@@
+
+    mma_warp_id = 4
+    tma_warp_id = 5
+    scheduler_warp_id = 6
+    unused_warp_id = 7
+    num_epilogue_warps = 4
+    epi_reg_count = 232
+    prod_reg_count = 24
+
+    warp_idx = cute.arch.warp_idx()
+    warp_idx = cute.arch.make_warp_uniform(warp_idx)
+    elect_one = nvvm.elect_sync()
+
+    tidx = cute.arch.thread_idx()[0]
+    bidx = cute.arch.block_idx()[0]
+    bidy = cute.arch.block_idx()[1]
+    bidz = cute.arch.block_idx()[2]
+    gridx = cute.arch.grid_dim()[0]
+    gridy = cute.arch.grid_dim()[1]
+
+    cluster_m = cluster_shape_mnk[0]
+    cluster_n = cluster_shape_mnk[1]
+    cluster_size = cluster_m * cluster_n * cluster_shape_mnk[2]
+
+    cta_rank_in_cluster = cute.arch.block_idx_in_cluster()
+    m_rank = cta_rank_in_cluster % cluster_m
+    n_rank = cta_rank_in_cluster // cluster_m
+
+    if warp_idx == mma_warp_id:
+        for _i in cutlass.range_constexpr(num_a_operands):
+            nvvm.prefetch_tensormap(tma_a_descs[_i].get_ptr())
+            nvvm.prefetch_tensormap(tma_sfa_descs[_i].get_ptr())
+        for _j in cutlass.range_constexpr(num_b_operands):
+            nvvm.prefetch_tensormap(tma_b_descs[_j].get_ptr())
+            nvvm.prefetch_tensormap(tma_sfb_descs[_j].get_ptr())
+
+    cluster_linear_init = bidx // cluster_m
+
+    a_pattern = 0
+    for n_idx in cutlass.range_constexpr(cluster_n):
+        a_pattern = a_pattern | (1 << (n_idx * cluster_m))
+    b_pattern = (1 << cluster_m) - 1
+
+    if cutlass.const_expr(multicast_a):
+        tma_mcast_mask_a = cutlass.Int16(a_pattern) << m_rank
+    else:
+        tma_mcast_mask_a = cutlass.Int16(1) << cta_rank_in_cluster
+    if cutlass.const_expr(multicast_b):
+        tma_mcast_mask_b = cutlass.Int16(b_pattern) << (n_rank * cluster_m)
+    else:
+        tma_mcast_mask_b = cutlass.Int16(1) << cta_rank_in_cluster
+
+    a_part_arrive = cutlass.Int16(a_pattern) << m_rank
+    b_part_arrive = cutlass.Int16(b_pattern) << (n_rank * cluster_m)
+    ab_empty_arrive_mask = a_part_arrive | b_part_arrive
+
+    _smem_sys_reserved = cutlass.Array(cutlass.Int8, 1024, space=cutlass.AddressSpace.smem, alignment=1)
+
+    ab_full_mbar_ptr = cutlass.Array(cutlass.Int64, ab_stages, space=cutlass.AddressSpace.smem)
+    ab_empty_mbar_ptr = cutlass.Array(cutlass.Int64, ab_stages, space=cutlass.AddressSpace.smem)
+    acc_empty_mbar_ptr = cutlass.Array(cutlass.Int64, acc_stages, space=cutlass.AddressSpace.smem)
+    acc_full_mbar_ptr = cutlass.Array(cutlass.Int64, acc_stages, space=cutlass.AddressSpace.smem)
+
+    tmem_dealloc_mbar_ptr = cutlass.Array(cutlass.Int64, 1, space=cutlass.AddressSpace.smem)
+    tmem_ptr_i32 = cutlass.Array(cutlass.Int32, 1, space=cutlass.AddressSpace.smem)
+
+    sched_storage = cutlass.Array(
+        cutlass.Int32,
+        SCHED_STAGES * SCHED_SLOT_WORDS,
+        space=cutlass.AddressSpace.smem,
+        alignment=16,
+    )
+    sched_full_mbar_ptr = cutlass.Array(cutlass.Int64, SCHED_STAGES, space=cutlass.AddressSpace.smem, alignment=8)
+    sched_empty_mbar_ptr = cutlass.Array(cutlass.Int64, SCHED_STAGES, space=cutlass.AddressSpace.smem, alignment=8)
+    tma_a_desc_smem_list = [
+        cutlass.Array(
+            cutlass.Int64,
+            TENSOR_MAP_QWORDS,
+            space=cutlass.AddressSpace.smem,
+            alignment=128,
+        )
+        for _ in range(num_a_operands)
+    ]
+
+    sA_elems = sA_packed_elems
+    sB_elems = sB_packed_elems
+    smem_a_list = [
+        cutlass.Array(
+            ab_dtype,
+            sA_elems * ab_stages,
+            space=cutlass.AddressSpace.smem,
+            alignment=1024,
+        )
+        for _ in range(num_a_operands)
+    ]
+    smem_b_list = [
+        cutlass.Array(
+            ab_dtype,
+            sB_elems * ab_stages,
+            space=cutlass.AddressSpace.smem,
+            alignment=1024,
+        )
+        for _ in range(num_b_operands)
+    ]
+    smem_sfa_list = [
+        cutlass.Array(
+            cutlass.Uint8,
+            sfa_smem_bytes * ab_stages,
+            space=cutlass.AddressSpace.smem,
+            alignment=1024,
+        )
+        for _ in range(num_a_operands)
+    ]
+    smem_sfb_list = [
+        cutlass.Array(
+            cutlass.Uint8,
+            sfb_smem_bytes * ab_stages,
+            space=cutlass.AddressSpace.smem,
+            alignment=1024,
+        )
+        for _ in range(num_b_operands)
+    ]
+
+    ab_empty_count = cluster_m + cluster_n - 1
+    sched_empty_count = 1 + 1 + num_epilogue_warps
+    if warp_idx == 0:
+        for i in range(ab_stages):
+            if elect_one:
+                nvvm.mbarrier_init(ab_full_mbar_ptr.subview(i), 1)
+            if elect_one:
+                nvvm.mbarrier_init(ab_empty_mbar_ptr.subview(i), ab_empty_count)
+        for i in range(acc_stages):
+            if elect_one:
+                nvvm.mbarrier_init(acc_full_mbar_ptr.subview(i), 1)
+            if elect_one:
+                nvvm.mbarrier_init(acc_empty_mbar_ptr.subview(i), num_epilogue_warps)
+        if cutlass.const_expr(use_acc_overlap):
+            if elect_one:
+                nvvm.mbarrier_init(tmem_dealloc_mbar_ptr, num_epilogue_warps)
+        for i in range(SCHED_STAGES):
+            if elect_one:
+                nvvm.mbarrier_init(sched_full_mbar_ptr.subview(i), 1)
+            if elect_one:
+                nvvm.mbarrier_init(sched_empty_mbar_ptr.subview(i), sched_empty_count)
+    nvvm.fence_mbarrier_init()
+
+    if cutlass.const_expr(cluster_shape_mnk[0] * cluster_shape_mnk[1] > 1):
+        nvvm.barrier_cluster_arrive_relaxed()
+        nvvm.barrier_cluster_wait()
+    else:
+        nvvm.barrier_cta_sync(0)
+
+    sA_bytes = sA_elems * (ab_dtype.width // 8)
+    sB_bytes = sB_elems * (ab_dtype.width // 8)
+
+    num_tma_copy_bytes = num_a_operands * (sA_bytes + sfa_smem_bytes) + num_b_operands * (sB_bytes + sfb_smem_bytes)
+
+    cols_per_acc_stage = cta_tile_mnk[1]
+    tmem_alloc_bar_count = (num_epilogue_warps + 1) * 32
+
+    # @@INJECT_TAP_PTRS@@
+
+    VEC_BYTES = vec_bytes_epi
+    vsize = (VEC_BYTES * 8) // cd_dtype.width
+
+    M = m
+    N = n
+    clusters_along_n = cute.ceil_div(cutlass.Int32(N), cgrp_tile_mnk[1])
+    num_k_tiles = cute.ceil_div(k, cta_tile_mnk[2])
+    first_token_arr = cutlass.make_array_view(first_token_offset)
+
+    if warp_idx == scheduler_warp_id:
+        nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
+        full_warp_mask = 0xFFFFFFFF
+        shfl_idx_clamp = 0x1F
+        shfl_up_clamp = 0
+        lane = cute.arch.lane_idx()
+        gemm_s = cutlass.Int32(M)
+        sched_stage = cutlass.Int32(0)
+        sched_empty_phase = cutlass.Int32(1)
+        linear_idx = cutlass.Int32(cluster_linear_init)
+        start_linear_idx = cutlass.Int32(0)
+        total_tiles = cutlass.Int32(0)
+        start_sf_block_m = cutlass.Int32(0)
+        total_sf_blocks_m = cutlass.Int32(0)
+        group_idx = cutlass.Int32(0)
+        is_tile_valid = cutlass.Int32(1)
+        cached_next_end = cutlass.Int32(0)
+        if lane + 1 < num_groups:
+            cached_next_end = cutlass.Int32(first_token_arr[lane + 1])
+        else:
+            cached_next_end = gemm_s
+        tile_lower_bound = nvvm.shfl_sync(full_warp_mask, cached_next_end, 1, shfl_up_clamp, nvvm.Shfl.UP)
+        cached_next_begin = cutlass.Int32(0)
+        if lane != 0:
+            cached_next_begin = tile_lower_bound
+
+        while is_tile_valid != 0:
+            group_begin = cached_next_begin
+            group_end = cached_next_end
+
+            if linear_idx >= start_linear_idx + total_tiles:
+                group_idx += lane
+                is_search_live = cutlass.Int32(1)
+                while is_search_live != 0:
+                    cached_group_begin = cached_next_begin
+                    cached_group_end = cached_next_end
+                    tile_start_idx = nvvm.shfl_sync(
+                        full_warp_mask,
+                        cached_next_end,
+                        31,
+                        shfl_idx_clamp,
+                        nvvm.Shfl.IDX,
+                    )
+                    next_end_group = group_idx + 32 + 1
+                    if next_end_group < num_groups:
+                        cached_next_end = cutlass.Int32(first_token_arr[next_end_group])
+                    else:
+                        cached_next_end = gemm_s
+                    tile_lower_bound = nvvm.shfl_sync(
+                        full_warp_mask,
+                        cached_next_end,
+                        1,
+                        shfl_up_clamp,
+                        nvvm.Shfl.UP,
+                    )
+                    if lane != 0:
+                        cached_next_begin = tile_lower_bound
+                    else:
+                        cached_next_begin = tile_start_idx
+
+                    group_m = cached_group_end - cached_group_begin
+                    total_tiles = cute.ceil_div(group_m, cgrp_tile_mnk[0]) * clusters_along_n
+                    total_sf_blocks_m = cute.ceil_div(group_m, 128)
+                    prefix_tiles = total_tiles
+                    prefix_sf = total_sf_blocks_m
+                    for delta in (1, 2, 4, 8, 16):
+                        prefix_delta = nvvm.shfl_sync(
+                            full_warp_mask,
+                            prefix_tiles,
+                            delta,
+                            shfl_up_clamp,
+                            nvvm.Shfl.UP,
+                        )
+                        prefix_sf_delta = nvvm.shfl_sync(
+                            full_warp_mask,
+                            prefix_sf,
+                            delta,
+                            shfl_up_clamp,
+                            nvvm.Shfl.UP,
+                        )
+                        if lane >= delta:
+                            prefix_tiles += prefix_delta
+                            prefix_sf += prefix_sf_delta
+                    start_linear_idx += prefix_tiles - total_tiles
+                    start_sf_block_m += prefix_sf - total_sf_blocks_m
+                    thread_succeed = nvvm.vote_sync(
+                        full_warp_mask,
+                        linear_idx < start_linear_idx + total_tiles,
+                        nvvm.VoteSync.BALLOT,
+                    )
+                    if thread_succeed != 0:
+                        winning_lane = cutlass.Int32(31) - cute.arch.bfind(cute.arch.brev(thread_succeed)).to(cutlass.Int32)
+                        group_idx = nvvm.shfl_sync(
+                            full_warp_mask,
+                            group_idx,
+                            winning_lane,
+                            shfl_idx_clamp,
+                            nvvm.Shfl.IDX,
+                        )
+                        start_linear_idx = nvvm.shfl_sync(
+                            full_warp_mask,
+                            start_linear_idx,
+                            winning_lane,
+                            shfl_idx_clamp,
+                            nvvm.Shfl.IDX,
+                        )
+                        total_tiles = nvvm.shfl_sync(
+                            full_warp_mask,
+                            total_tiles,
+                            winning_lane,
+                            shfl_idx_clamp,
+                            nvvm.Shfl.IDX,
+                        )
+                        start_sf_block_m = nvvm.shfl_sync(
+                            full_warp_mask,
+                            start_sf_block_m,
+                            winning_lane,
+                            shfl_idx_clamp,
+                            nvvm.Shfl.IDX,
+                        )
+                        tile_start_idx = nvvm.shfl_sync(
+                            full_warp_mask,
+                            cached_group_begin,
+                            winning_lane,
+                            shfl_idx_clamp,
+                            nvvm.Shfl.IDX,
+                        )
+                        group_end_idx = group_idx + lane + 1
+                        if group_end_idx < num_groups:
+                            cached_next_end = cutlass.Int32(first_token_arr[group_end_idx])
+                        else:
+                            cached_next_end = gemm_s
+                        tile_lower_bound = nvvm.shfl_sync(
+                            full_warp_mask,
+                            cached_next_end,
+                            1,
+                            shfl_up_clamp,
+                            nvvm.Shfl.UP,
+                        )
+                        if lane != 0:
+                            cached_next_begin = tile_lower_bound
+                        else:
+                            cached_next_begin = tile_start_idx
+                        group_begin = cached_next_begin
+                        group_end = cached_next_end
+                        is_search_live = cutlass.Int32(0)
+                    else:
+                        group_idx += 32
+                        first_lane_group = nvvm.shfl_sync(
+                            full_warp_mask,
+                            group_idx,
+                            0,
+                            shfl_idx_clamp,
+                            nvvm.Shfl.IDX,
+                        )
+                        if first_lane_group >= num_groups:
+                            is_tile_valid = cutlass.Int32(0)
+                            is_search_live = cutlass.Int32(0)
+                        else:
+                            next_start_linear_idx = start_linear_idx + total_tiles
+                            start_linear_idx = nvvm.shfl_sync(
+                                full_warp_mask,
+                                next_start_linear_idx,
+                                31,
+                                shfl_idx_clamp,
+                                nvvm.Shfl.IDX,
+                            )
+                            next_start_sf = start_sf_block_m + total_sf_blocks_m
+                            start_sf_block_m = nvvm.shfl_sync(
+                                full_warp_mask,
+                                next_start_sf,
+                                31,
+                                shfl_idx_clamp,
+                                nvvm.Shfl.IDX,
+                            )
+
+            coord_expert = cutlass.Int32(0)
+            cluster_tile_m = cutlass.Int32(0)
+            coord_n = cutlass.Int32(0)
+            if is_tile_valid != 0:
+                local_linear_idx = linear_idx - start_linear_idx
+                group_nt_m = total_tiles // clusters_along_n
+                cluster_tile_m, coord_n = _moe_swizzle_tile(
+                    local_linear_idx,
+                    group_nt_m,
+                    clusters_along_n,
+                    _moe_auto_swizzle_w(group_nt_m * cgrp_tile_mnk[0], N, k, clusters_along_n),
+                )
+                coord_expert = group_idx % num_experts
+                linear_idx += grid_num_clusters
+
+            while not nvvm.mbarrier_try_wait_parity(
+                sched_empty_mbar_ptr.subview(sched_stage),
+                sched_empty_phase,
+                time_limit=10_000_000,
+            ):
+                pass
+            if lane == 0:
+                slot = sched_storage.subview(sched_stage * SCHED_SLOT_WORDS)
+                (slot.subview(0)).store(coord_expert)
+                (slot.subview(1)).store(cluster_tile_m)
+                (slot.subview(2)).store(coord_n)
+                (slot.subview(3)).store(is_tile_valid)
+                (slot.subview(4)).store(group_begin)
+                (slot.subview(5)).store(group_end)
+                (slot.subview(6)).store(start_sf_block_m)
+                (slot.subview(7)).store(group_idx)
+                nvvm.mbarrier_arrive(sched_full_mbar_ptr.subview(sched_stage))
+
+            sched_stage += 1
+            if sched_stage == SCHED_STAGES:
+                sched_stage = cutlass.Int32(0)
+                sched_empty_phase = sched_empty_phase ^ 1
+
+    if warp_idx == tma_warp_id:
+        nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
+        if cutlass.const_expr(USE_PDL):
+            if elect_one:
+                nvvm.griddepcontrol("wait")
+        ab_empty_phase_bit = cutlass.Int32(1)
+        ab_iter = cutlass.Int32(0)
+        sched_stage = cutlass.Int32(0)
+        sched_full_phase = cutlass.Int32(0)
+        is_valid = cutlass.Int32(1)
+
+        lane = tidx % 32
+        block_linear = bidx + bidy * gridx
+        cta_desc_base_list = [a_tma_workspace.iterator.raw_ptr() + (block_linear * num_a_operands + _ai) * TENSOR_MAP_QWORDS for _ai in range(num_a_operands)]
+        a_desc_tma_ptr_list = [
+            cute.make_ptr(
+                cutlass.Int64,
+                cta_desc_base_list[_ai].toint(),
+                mem_space=cute.AddressSpace.generic,
+            )
+            for _ai in range(num_a_operands)
+        ]
+        previous_group_begin = cutlass.Int32(-1)
+        if elect_one:
+            for _ai in cutlass.range_constexpr(num_a_operands):
+                _copy_tensormap_to_workspace(tma_a_descs[_ai].get_ptr(), tma_a_desc_smem_list[_ai])
+        nvvm.bar_warp_sync(0xFFFFFFFF)
+
+        while is_valid != 0:
+            while not nvvm.mbarrier_try_wait_parity(
+                sched_full_mbar_ptr.subview(sched_stage),
+                sched_full_phase,
+                time_limit=10_000_000,
+            ):
+                pass
+            slot = sched_storage.subview(sched_stage * SCHED_SLOT_WORDS)
+            coord_expert = (slot.subview(0)).load()
+            tile_m = (slot.subview(1)).load()
+            tile_n = (slot.subview(2)).load()
+            is_valid = (slot.subview(3)).load()
+            group_begin = (slot.subview(4)).load()
+            group_end = (slot.subview(5)).load()
+            start_sf_block_m = (slot.subview(6)).load()
+            if elect_one:
+                nvvm.mbarrier_arrive(sched_empty_mbar_ptr.subview(sched_stage))
+            sched_stage += 1
+            if sched_stage == SCHED_STAGES:
+                sched_stage = cutlass.Int32(0)
+                sched_full_phase = sched_full_phase ^ 1
+
+            if is_valid != 0:
+                coord_m_group = tile_m * cgrp_tile_mnk[0] + m_rank * cta_tile_mnk[0]
+                coord_n_per_cta = tile_n * cgrp_tile_mnk[1] + n_rank * cta_tile_mnk[1]
+                sfa_m_block = start_sf_block_m + coord_m_group // 128
+                sfb_n_block = coord_n_per_cta // 128
+
+                if group_begin != previous_group_begin:
+                    previous_group_begin = group_begin
+                    for _ai in cutlass.range_constexpr(num_a_operands):
+                        _fence_tensormap_acquire(a_desc_tma_ptr_list[_ai])
+                    for _ai in cutlass.range_constexpr(num_a_operands):
+                        if elect_one:
+                            row_base = mA_list[_ai].iterator.raw_ptr().toint() + ((group_begin * a_stride_m_list[_ai] * ab_dtype.width) >> 3)
+                            _replace_tensormap_global_address(tma_a_desc_smem_list[_ai], row_base)
+                            _replace_tensormap_global_dim_1(tma_a_desc_smem_list[_ai], group_end - group_begin)
+                        nvvm.bar_warp_sync(0xFFFFFFFF)
+                        if lane < TENSOR_MAP_QWORDS:
+                            (cta_desc_base_list[_ai] + lane).store((tma_a_desc_smem_list[_ai].subview(lane)).load())
+                        nvvm.bar_warp_sync(0xFFFFFFFF)
+                        _fence_tensormap_release()
+
+                for k_tile_idx in range(num_k_tiles):
+                    stage = ab_iter % ab_stages
+                    if stage == 0 and ab_iter != 0:
+                        ab_empty_phase_bit = ab_empty_phase_bit ^ 1
+
+                    while not nvvm.mbarrier_try_wait_parity(
+                        ab_empty_mbar_ptr.subview(stage),
+                        ab_empty_phase_bit,
+                        time_limit=10_000_000,
+                    ):
+                        pass
+
+                    coord_k = k_tile_idx * cta_tile_mnk[2]
+                    coord_sf_k = k_tile_idx * sf_tma_box_k
+                    if elect_one:
+                        nvvm.mbarrier_arrive_expect_tx(ab_full_mbar_ptr.subview(stage), num_tma_copy_bytes)
+                    a_issue = (not multicast_a) or (n_rank == 0)
+                    b_issue = (not multicast_b) or (m_rank == 0)
+                    if a_issue:
+                        for _ai in cutlass.range_constexpr(num_a_operands):
+                            if elect_one:
+                                nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                    smem_a_list[_ai].subview(sA_elems * stage),
+                                    a_desc_tma_ptr_list[_ai],
+                                    (coord_k, coord_m_group, cutlass.Int32(0)),
+                                    ab_full_mbar_ptr.subview(stage),
+                                    [],
+                                    multicast_mask=tma_mcast_mask_a,
+                                    group=nvvm.CTAGroup.CTA_1,
+                                )
+                        for _ai in cutlass.range_constexpr(num_a_operands):
+                            if elect_one:
+                                nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                    smem_sfa_list[_ai].subview(sfa_smem_bytes * stage),
+                                    tma_sfa_descs[_ai].get_ptr(),
+                                    (0, coord_sf_k, sfa_m_block, cutlass.Int32(0)),
+                                    ab_full_mbar_ptr.subview(stage),
+                                    [],
+                                    multicast_mask=tma_mcast_mask_a,
+                                    group=nvvm.CTAGroup.CTA_1,
+                                )
+                    if b_issue:
+                        for _bj in cutlass.range_constexpr(num_b_operands):
+                            sB_stage = smem_b_list[_bj].subview(sB_elems * stage)
+                            if cutlass.const_expr(b_is_n_major):
+                                for n_group in cutlass.range_constexpr(cta_tile_mnk[1] // b_tma_group_elems):
+                                    if elect_one:
+                                        nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                            sB_stage.subview(n_group * b_tma_group_elems * cta_tile_mnk[2]),
+                                            tma_b_descs[_bj].get_ptr(),
+                                            (
+                                                coord_n_per_cta + n_group * b_tma_group_elems,
+                                                coord_k,
+                                                coord_expert,
+                                            ),
+                                            ab_full_mbar_ptr.subview(stage),
+                                            [],
+                                            multicast_mask=tma_mcast_mask_b,
+                                            group=nvvm.CTAGroup.CTA_1,
+                                        )
+                            else:
+                                if elect_one:
+                                    nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                        sB_stage,
+                                        tma_b_descs[_bj].get_ptr(),
+                                        (coord_k, coord_n_per_cta, coord_expert),
+                                        ab_full_mbar_ptr.subview(stage),
+                                        [],
+                                        multicast_mask=tma_mcast_mask_b,
+                                        group=nvvm.CTAGroup.CTA_1,
+                                    )
+                        for _bj in cutlass.range_constexpr(num_b_operands):
+                            if elect_one:
+                                nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                    smem_sfb_list[_bj].subview(sfb_smem_bytes * stage),
+                                    tma_sfb_descs[_bj].get_ptr(),
+                                    (0, coord_sf_k, sfb_n_block, coord_expert),
+                                    ab_full_mbar_ptr.subview(stage),
+                                    [],
+                                    multicast_mask=tma_mcast_mask_b,
+                                    group=nvvm.CTAGroup.CTA_1,
+                                )
+                    ab_iter += 1
+
+        tail_stage = ab_iter % ab_stages
+        tail_phase = ab_empty_phase_bit
+        if tail_stage == 0 and ab_iter != 0:
+            tail_phase = tail_phase ^ 1
+        for _ in range(ab_stages - 1):
+            tail_stage = tail_stage + 1
+            if tail_stage == ab_stages:
+                tail_stage = cutlass.Int32(0)
+                tail_phase = tail_phase ^ 1
+        if elect_one:
+            while not nvvm.mbarrier_try_wait_parity(ab_empty_mbar_ptr.subview(tail_stage), tail_phase, time_limit=10_000_000):
+                pass
+
+    if warp_idx == mma_warp_id:
+        nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
+        nvvm.tcgen05_alloc(
+            tmem_ptr_i32,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_1,
+        )
+        nvvm.bar_warp_sync(0xFFFFFFFF)
+        nvvm.barrier_cta_arrive(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
+        tmem_raw_addr = tmem_ptr_i32.load()
+        base_col_id_root = tmem_raw_addr & 0xFFFF
+        base_row_id = tmem_raw_addr >> 16
+        ab_full_phase_bit = cutlass.Int32(0)
+        ab_iter = cutlass.Int32(0)
+        acc_empty_phase_bit = cutlass.Int32(1)
+        tile_iter = cutlass.Int32(0)
+        is_valid = cutlass.Int32(1)
+        sched_stage = cutlass.Int32(0)
+        sched_full_phase = cutlass.Int32(0)
+        acc_stage = cutlass.Int32(0)
+
+        # fp4 packs its K-mode into the OMMA descriptor's 2-bit split field; fp8
+        # keeps the MX descriptor's 1-bit one. Both are built once, outside the
+        # loops — the fields depend only on j (the scale id within a word).
+        if cutlass.const_expr(idesc_is_omma):
+            idesc_by_j = [
+                cutlass.experimental.primitives.Tcgen05MxOmmaInstrDesc.build(
+                    a_dtype=idesc_a_dtype,
+                    b_dtype=idesc_b_dtype,
+                    scale_format=sf_scale_format,
+                    n_dim=mma_n_dim,
+                    m_dim=mma_m_dim,
+                    a_major=mma_a_major,
+                    b_major=mma_b_major,
+                    a_sf_id=j * sf_scales_per_inst,
+                    b_sf_id=j * sf_scales_per_inst,
+                    k_dim=mma_k_dim_mode,
+                )
+                for j in range(sf_insts_per_atom)
+            ]
+        else:
+            idesc_by_j = [
+                cutlass.experimental.primitives.Tcgen05MxInstrDesc.build(
+                    a_dtype=idesc_a_dtype,
+                    b_dtype=idesc_b_dtype,
+                    scale_format=sf_scale_format,
+                    n_dim=mma_n_dim,
+                    m_dim=mma_m_dim,
+                    a_major=mma_a_major,
+                    b_major=mma_b_major,
+                    a_sf_id=j * sf_scales_per_inst,
+                    b_sf_id=j * sf_scales_per_inst,
+                    k_dim=mma_k_dim_mode,
+                )
+                for j in range(sf_insts_per_atom)
+            ]
+
+        sfa_tmem_bases = [(base_row_id << 16) | (base_col_id_root + sfa_col_bases[i]) for i in range(num_a_operands)]
+        sfb_tmem_bases = [(base_row_id << 16) | (base_col_id_root + sfb_col_bases[j]) for j in range(num_b_operands)]
+        s2t_shape, s2t_multicast = nvvm.S2TCopyMode.S2T_32x128b_WARPX4
+        sfa_scale_ptrs = [nvvm.make_tmem_ptr(b, cutlass.Float32) for b in sfa_tmem_bases]
+        sfb_scale_ptrs = [nvvm.make_tmem_ptr(b, cutlass.Float32) for b in sfb_tmem_bases]
+        # utccp destination per (MN-block, atom within the scale word). SFB is
+        # atom-MAJOR across the N-blocks because ONE instruction walks all of
+        # them; SFA is block-major because one instruction covers exactly one
+        # 128-row block, so that word has to be contiguous. Both collapse to
+        # the same addresses at a single block, and to sm100's layout at
+        # word_atoms == 1.
+        sfa_dst_ptrs = [
+            [
+                [nvvm.make_tmem_ptr(sfa_tmem_bases[i] + m * registers_per_block + a * registers_per_atom, cutlass.Float32) for a in range(word_atoms)]
+                for m in range(num_blocks_m)
+            ]
+            for i in range(num_a_operands)
+        ]
+        sfb_dst_ptrs = [
+            [
+                [nvvm.make_tmem_ptr(sfb_tmem_bases[j] + (a * num_blocks_n + m) * registers_per_atom, cutlass.Float32) for a in range(word_atoms)]
+                for m in range(num_blocks_n)
+            ]
+            for j in range(num_b_operands)
+        ]
+        while is_valid != 0:
+            while not nvvm.mbarrier_try_wait_parity(
+                sched_full_mbar_ptr.subview(sched_stage),
+                sched_full_phase,
+                time_limit=10_000_000,
+            ):
+                pass
+            is_valid = (sched_storage.subview(sched_stage * SCHED_SLOT_WORDS).subview(3)).load()
+            if elect_one:
+                nvvm.mbarrier_arrive(sched_empty_mbar_ptr.subview(sched_stage))
+            sched_stage += 1
+            if sched_stage == SCHED_STAGES:
+                sched_stage = cutlass.Int32(0)
+                sched_full_phase = sched_full_phase ^ 1
+
+            if is_valid != 0:
+                acc_stage = tile_iter % acc_stages
+                if acc_stage == 0 and tile_iter != 0:
+                    acc_empty_phase_bit = acc_empty_phase_bit ^ 1
+
+                while not nvvm.mbarrier_try_wait_parity(
+                    acc_empty_mbar_ptr.subview(acc_stage),
+                    acc_empty_phase_bit,
+                    time_limit=10_000_000,
+                ):
+                    pass
+
+                if cutlass.const_expr(use_acc_overlap):
+                    acc_base_col = base_col_id_root + (tile_iter % 2) * acc_stage_stride
+                else:
+                    acc_base_col = base_col_id_root + acc_stage * acc_region_cols
+                # One accumulator per (gemm, M block); M block mi sits
+                # epi_cols_per_mma_m columns further into its GEMM's region and
+                # reads SF word block mi (SF words are one per 128 rows).
+                acc_tmem_ptrs = [
+                    [
+                        nvvm.make_tmem_ptr(
+                            (base_row_id << 16) | (acc_base_col + g * acc_gemm_stride + mi * epi_cols_per_mma_m),
+                            cutlass.Float32,
+                        )
+                        for mi in range(num_mma_m)
+                    ]
+                    for g in range(num_gemms)
+                ]
+
+                scale_d = cutlass.Boolean(False)
+                for k_tile_idx in range(num_k_tiles):
+                    stage = ab_iter % ab_stages
+                    if stage == 0 and ab_iter != 0:
+                        ab_full_phase_bit = ab_full_phase_bit ^ 1
+
+                    while not nvvm.mbarrier_try_wait_parity(
+                        ab_full_mbar_ptr.subview(stage),
+                        ab_full_phase_bit,
+                        time_limit=10_000_000,
+                    ):
+                        pass
+
+                    desc_a_bases = [
+                        cutlass.experimental.primitives.Tcgen05SmemDesc.build(
+                            start_address=smem_a_list[i].subview(sA_elems * stage),
+                            leading_byte_offset=a_smem_desc_leading_byte_offset,
+                            stride_byte_offset=a_smem_desc_stride_byte_offset,
+                            layout=ab_smem_swizzle,
+                        )
+                        for i in range(num_a_operands)
+                    ]
+                    desc_b_bases = [
+                        cutlass.experimental.primitives.Tcgen05SmemDesc.build(
+                            start_address=smem_b_list[j].subview(sB_elems * stage),
+                            leading_byte_offset=b_smem_desc_leading_byte_offset,
+                            stride_byte_offset=b_smem_desc_stride_byte_offset,
+                            layout=ab_smem_swizzle,
+                        )
+                        for j in range(num_b_operands)
+                    ]
+                    desc_sfa_bases = [
+                        cutlass.experimental.primitives.Tcgen05SmemDesc.build(
+                            start_address=smem_sfa_list[i].subview(sfa_smem_bytes * stage),
+                            leading_byte_offset=16,
+                            stride_byte_offset=128,
+                            layout=cutlass.experimental.primitives.Tcgen05SmemSwizzle.NONE,
+                        )
+                        for i in range(num_a_operands)
+                    ]
+                    desc_sfb_bases = [
+                        cutlass.experimental.primitives.Tcgen05SmemDesc.build(
+                            start_address=smem_sfb_list[j].subview(sfb_smem_bytes * stage),
+                            leading_byte_offset=16,
+                            stride_byte_offset=128,
+                            layout=cutlass.experimental.primitives.Tcgen05SmemSwizzle.NONE,
+                        )
+                        for j in range(num_b_operands)
+                    ]
+
+                    # One SF word per group of MMAs, refreshed right before they
+                    # read it. A word spans word_atoms consecutive K-atoms in SMEM.
+                    for atom_r in cutlass.range_constexpr(num_sf_atoms):
+                        for _ai in cutlass.range_constexpr(num_a_operands):
+                            for _m in cutlass.range_constexpr(num_blocks_m):
+                                for _a in cutlass.range_constexpr(word_atoms):
+                                    if elect_one:
+                                        nvvm.tcgen05_cp(
+                                            s2t_shape,
+                                            sfa_dst_ptrs[_ai][_m][_a],
+                                            desc_sfa_bases[_ai] + (sf_atom_desc_stride * (atom_r * word_atoms + _a) + sf_block_desc_stride * _m),
+                                            group=nvvm.CTAGroup.CTA_1,
+                                            multicast=s2t_multicast,
+                                        )
+                        for _bj in cutlass.range_constexpr(num_b_operands):
+                            for _m in cutlass.range_constexpr(num_blocks_n):
+                                for _a in cutlass.range_constexpr(word_atoms):
+                                    if elect_one:
+                                        nvvm.tcgen05_cp(
+                                            s2t_shape,
+                                            sfb_dst_ptrs[_bj][_m][_a],
+                                            desc_sfb_bases[_bj] + (sf_atom_desc_stride * (atom_r * word_atoms + _a) + sf_block_desc_stride * _m),
+                                            group=nvvm.CTAGroup.CTA_1,
+                                            multicast=s2t_multicast,
+                                        )
+                        for j in cutlass.range_constexpr(sf_insts_per_atom):
+                            k_block_idx = atom_r * sf_insts_per_atom + j
+                            idesc_k = idesc_by_j[j]
+                            for g in cutlass.range_constexpr(num_gemms):
+                                _ai = gemm_a_idx[g]
+                                _bj = gemm_b_idx[g]
+                                desc_a_k = desc_a_bases[_ai].advance_start_address(a_smem_k_step_bytes * k_block_idx)
+                                desc_b = desc_b_bases[_bj].advance_start_address(b_smem_k_step_bytes * k_block_idx)
+                                for mi in cutlass.range_constexpr(num_mma_m):
+                                    # The M sub-block offset is a whole SMEM swizzle atom, so
+                                    # the descriptor's swizzle phase is preserved. B and its SF
+                                    # are shared; A's SF word block follows the M block.
+                                    desc_a = desc_a_k.advance_start_address(a_smem_m_step_bytes * mi)
+                                    if elect_one:
+                                        nvvm.tcgen05_mma_block_scale(
+                                            mma_block_scale_kind,
+                                            nvvm.CTAGroup.CTA_1,
+                                            acc_tmem_ptrs[g][mi],
+                                            desc_a,
+                                            desc_b,
+                                            idesc_k,
+                                            enable_input_d=scale_d,
+                                            scale_a=sfa_dst_ptrs[_ai][mi][0],
+                                            scale_b=sfb_scale_ptrs[_bj],
+                                            scale_vec_size=scale_vec_size,
+                                            b_collector_op=_b_collector_op(mi),
+                                        )
+                            # Every accumulator sees scale_d=False on exactly the first
+                            # k_block of the tile, so the flip stays outside mi.
+                            scale_d = cutlass.Boolean(True)
+
+                    if elect_one:
+                        nvvm.tcgen05_commit(
+                            ab_empty_mbar_ptr.subview(stage),
+                            multicast_mask=ab_empty_arrive_mask,
+                            group=nvvm.CTAGroup.CTA_1,
+                        )
+                    ab_iter += 1
+
+                if elect_one:
+                    nvvm.tcgen05_commit(
+                        acc_full_mbar_ptr.subview(acc_stage),
+                        group=nvvm.CTAGroup.CTA_1,
+                    )
+                tile_iter += 1
+
+        if cutlass.const_expr(USE_PDL):
+            if elect_one:
+                nvvm.griddepcontrol("launch_dependents")
+
+        nvvm.tcgen05_relinquish_alloc_permit(group=nvvm.CTAGroup.CTA_1)
+        if tile_iter != 0:
+            tail_stage = acc_stage
+            tail_phase = acc_empty_phase_bit
+            if elect_one:
+                for _ in range(acc_stages):
+                    tail_stage = tail_stage + 1
+                    if tail_stage == acc_stages:
+                        tail_stage = cutlass.Int32(0)
+                        tail_phase = tail_phase ^ 1
+                    while not nvvm.mbarrier_try_wait_parity(
+                        acc_empty_mbar_ptr.subview(tail_stage),
+                        tail_phase,
+                        time_limit=10_000_000,
+                    ):
+                        pass
+                if cutlass.const_expr(use_acc_overlap):
+                    while not nvvm.mbarrier_try_wait_parity(tmem_dealloc_mbar_ptr, 0, time_limit=10_000_000):
+                        pass
+
+        nvvm.bar_warp_sync(0xFFFFFFFF)
+        alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
+        nvvm.tcgen05_dealloc(
+            alloc_ptr,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_1,
+        )
+
+    if warp_idx < num_epilogue_warps:
+        nvvm.setmaxregister(epi_reg_count, nvvm.SetMaxRegisterAction.INCREASE)
+        nvvm.barrier_cta_sync(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
+        tmem_raw_addr = tmem_ptr_i32.load()
+        base_col_id_root = tmem_raw_addr & 0xFFFF
+        base_row_id = tmem_raw_addr >> 16
+
+        if cutlass.const_expr(USE_PDL):
+            nvvm.griddepcontrol("wait")
+
+        tile_iter = cutlass.Int32(0)
+        acc_full_phase_bit = cutlass.Int32(0)
+        is_valid = cutlass.Int32(1)
+        sched_stage = cutlass.Int32(0)
+        sched_full_phase = cutlass.Int32(0)
+
+        if cutlass.const_expr(mma_inst_shape_mnk[0] == 64):
+            row_id_with_warp_offset = base_row_id
+        else:
+            row_id_with_warp_offset = base_row_id + warp_idx * 32
+
+        # One M block's accumulator columns are contiguous.
+        subtile_cnt = cute.ceil_div(epi_cols_per_mma_m, 32)
+        t2r_inst_repx = epi_tile_mn[1]
+        if cutlass.const_expr(mma_inst_shape_mnk[0] == 64):
+            shape = nvvm.Tcgen05LdStShape.SHAPE_16X32BX2
+            ld_half_off = 0
+        else:
+            shape = nvvm.Tcgen05LdStShape.SHAPE_32X32B
+            ld_half_off = None
+        lane = tidx % 32
+
+        while is_valid != 0:
+            while not nvvm.mbarrier_try_wait_parity(
+                sched_full_mbar_ptr.subview(sched_stage),
+                sched_full_phase,
+                time_limit=10_000_000,
+            ):
+                pass
+            _slot = sched_storage.subview(sched_stage * SCHED_SLOT_WORDS)
+            tile_m = (_slot.subview(1)).load()
+            tile_n = (_slot.subview(2)).load()
+            is_valid = (_slot.subview(3)).load()
+            group_begin = (_slot.subview(4)).load()
+            group_end = (_slot.subview(5)).load()
+            group_idx = (_slot.subview(7)).load()
+            if elect_one:
+                nvvm.mbarrier_arrive(sched_empty_mbar_ptr.subview(sched_stage))
+            sched_stage += 1
+            if sched_stage == SCHED_STAGES:
+                sched_stage = cutlass.Int32(0)
+                sched_full_phase = sched_full_phase ^ 1
+
+            if is_valid != 0:
+                coord_m_tile = group_begin + tile_m * cgrp_tile_mnk[0] + m_rank * cta_tile_mnk[0]
+                coord_n = tile_n * cgrp_tile_mnk[1] + n_rank * cta_tile_mnk[1]
+
+                acc_stage = tile_iter % acc_stages
+                if acc_stage == 0 and tile_iter != 0:
+                    acc_full_phase_bit = acc_full_phase_bit ^ 1
+
+                while not nvvm.mbarrier_try_wait_parity(
+                    acc_full_mbar_ptr.subview(acc_stage),
+                    acc_full_phase_bit,
+                    time_limit=10_000_000,
+                ):
+                    pass
+
+                if cutlass.const_expr(use_acc_overlap):
+                    acc_buf_parity = tile_iter % 2
+                    acc_base_col = base_col_id_root + acc_buf_parity * acc_stage_stride
+                else:
+                    acc_buf_parity = cutlass.Int32(0)
+                    acc_base_col = base_col_id_root + acc_stage * acc_region_cols
+                # One pass per MMA-M block over its own column region.
+                for mi in cutlass.range_constexpr(num_mma_m):
+                    coord_m = coord_m_tile + mi * mma_inst_shape_mnk[0]
+                    mi_col_base = acc_base_col + mi * epi_cols_per_mma_m
+                    tmem_col_addr_gemms = [(row_id_with_warp_offset << 16) | (mi_col_base + g * acc_gemm_stride) for g in range(num_gemms)]
+
+                    if cutlass.const_expr(mma_inst_shape_mnk[0] == 64):
+                        row = coord_m + warp_idx * 16 + lane
+                        row_active = lane < 16
+                    else:
+                        row = coord_m + tidx
+                        row_active = True
+
+                    # @@INJECT_AUX_VIEWS@@
+
+                    for subtile_idx in cutlass.range_constexpr(subtile_cnt):
+                        if cutlass.const_expr(use_acc_overlap):
+                            _sub = subtile_idx + (1 - acc_buf_parity) * (subtile_cnt - 1 - 2 * subtile_idx)
+                            subtile_col_offset = _sub * 32
+                        else:
+                            subtile_col_offset = subtile_idx * 32
+                        c_rmem_vecs = []
+                        for g in cutlass.range_constexpr(num_gemms):
+                            tmem = cutlass.inttoptr(
+                                tmem_col_addr_gemms[g] + subtile_col_offset,
+                                6,
+                                cutlass.Float32,
+                            )
+                            c_rmem_vecs.append(nvvm.tcgen05_ld(shape, tmem, num=t2r_inst_repx, offset=ld_half_off))
+                        c_rmem_vec = c_rmem_vecs[0]
+
+                        if cutlass.const_expr(use_acc_overlap and mi == num_mma_m - 1 and subtile_idx == acc_overlap_subtiles - 1):
+                            nvvm.tcgen05_wait(kind=nvvm.Tcgen05Wait.LOAD)
+                            nvvm.tcgen05_fence(nvvm.Tcgen05Fence.BEFORE_THREAD_SYNC)
+                            if elect_one:
+                                nvvm.mbarrier_arrive(acc_empty_mbar_ptr.subview(acc_stage))
+
+                        col = coord_n + subtile_col_offset
+
+                        # @@STG_ONLY:BEGIN@@
+                        if row_active and row < group_end:
+                            for j in cutlass.range_constexpr(t2r_inst_repx // vsize):
+                                col_j = col + j * vsize
+                                if col_j + vsize <= N:
+                                    vec_f32 = c_rmem_vec[j * vsize : (j + 1) * vsize]
+
+                                    # @@INJECT_STG_VEC_BINDINGS@@
+
+                                    # @@INJECT_EPILOGUE@@
+                        # @@STG_ONLY:END@@
+
+                if cutlass.const_expr(not use_acc_overlap):
+                    nvvm.tcgen05_wait(kind=nvvm.Tcgen05Wait.LOAD)
+                    nvvm.tcgen05_fence(nvvm.Tcgen05Fence.BEFORE_THREAD_SYNC)
+                    if elect_one:
+                        nvvm.mbarrier_arrive(acc_empty_mbar_ptr.subview(acc_stage))
+                tile_iter += 1
+
+        if cutlass.const_expr(use_acc_overlap):
+            nvvm.tcgen05_wait(kind=nvvm.Tcgen05Wait.LOAD)
+            nvvm.tcgen05_fence(nvvm.Tcgen05Fence.BEFORE_THREAD_SYNC)
+            if elect_one:
+                nvvm.mbarrier_arrive(tmem_dealloc_mbar_ptr)
+
+    if warp_idx == unused_warp_id:
+        nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
+
+
+@cute.jit
+def _host(
+    problem_size: tuple,
+    first_token_offset: cute.Tensor,
+    a_tma_workspace: cute.Tensor,
+    # @@INJECT_HOST_AB_PARAMS@@
+    # @@INJECT_HOST_TAP_PARAMS@@
+    # @@INJECT_HOST_AUX_PARAMS@@
+    stream: _cuda.CUstream,
+) -> None:
+    # @@INJECT_HOST_AB_LISTS@@
+    m = problem_size[0]
+    n = problem_size[1]
+    k_sym = problem_size[2]
+    num_experts = problem_size[3]
+    num_groups = problem_size[4]
+    _stride_idx = 5
+    _a_stride_sets = []
+    for _ in cutlass.range_constexpr(num_a_operands):
+        _a_stride_sets.append(
+            (
+                problem_size[_stride_idx],
+                problem_size[_stride_idx + 1],
+                problem_size[_stride_idx + 2],
+            )
+        )
+        _stride_idx += 3
+    _b_stride_sets = []
+    for _ in cutlass.range_constexpr(num_b_operands):
+        _b_stride_sets.append(
+            (
+                problem_size[_stride_idx],
+                problem_size[_stride_idx + 1],
+                problem_size[_stride_idx + 2],
+            )
+        )
+        _stride_idx += 3
+    # @@INJECT_HOST_REDUCTION_STRIDES@@
+
+    tma_a_desc_list = []
+    for _a_idx, _a_op in enumerate(_a_operands):
+        a_stride_m, a_stride_k, a_stride_l = _a_stride_sets[_a_idx]
+        tma_a_desc_list.append(
+            _tma.create_tensor_map_tiled(
+                global_address=_a_op.iterator.toint(),
+                dtype=ab_tma_desc_dtype,
+                global_dims=[k_sym, m, 1],
+                global_strides=[
+                    a_stride_m * ab_dtype.width // 128,
+                    a_stride_l * ab_dtype.width // 128,
+                ],
+                box_dims=[cta_tile_mnk[2], cta_tile_mnk[0], 1],
+                swizzle=ab_tma_swizzle,
+                tma_format=ab_tma_format,
+            )
+        )
+    tma_b_desc_list = []
+    for _b_idx, _b_op in enumerate(_b_operands):
+        b_stride_n, b_stride_k, b_stride_l = _b_stride_sets[_b_idx]
+        if cutlass.const_expr(b_is_n_major):
+            tma_b_desc_list.append(
+                _tma.create_tensor_map_tiled(
+                    global_address=_b_op.iterator.toint(),
+                    dtype=ab_tma_desc_dtype,
+                    global_dims=[n, k_sym, num_experts],
+                    global_strides=[
+                        b_stride_k * ab_dtype.width // 128,
+                        b_stride_l * ab_dtype.width // 128,
+                    ],
+                    box_dims=[b_tma_group_elems, cta_tile_mnk[2], 1],
+                    swizzle=ab_tma_swizzle,
+                    tma_format=ab_tma_format,
+                )
+            )
+        else:
+            tma_b_desc_list.append(
+                _tma.create_tensor_map_tiled(
+                    global_address=_b_op.iterator.toint(),
+                    dtype=ab_tma_desc_dtype,
+                    global_dims=[k_sym, n, num_experts],
+                    global_strides=[
+                        b_stride_n * ab_dtype.width // 128,
+                        b_stride_l * ab_dtype.width // 128,
+                    ],
+                    box_dims=[cta_tile_mnk[2], cta_tile_mnk[1], 1],
+                    swizzle=ab_tma_swizzle,
+                    tma_format=ab_tma_format,
+                )
+            )
+    rest_k = ((k_sym // block_size) + 3) // 4
+    rest_m = (m + 127) // 128 + num_groups
+    rest_n = (n + 127) // 128
+    tma_sfa_desc_list = []
+    for _sfa_op in _sfa_operands:
+        sfa_fp16_tensor = cute.make_tensor(
+            cute.recast_ptr(_sfa_op.iterator, dtype=cutlass.Float16),
+            cute.make_layout(
+                (256, rest_k, rest_m, 1),
+                stride=(
+                    1,
+                    256,
+                    cute.assume(256 * rest_k, 8),
+                    cute.assume(256 * rest_k * rest_m, 8),
+                ),
+            ),
+        )
+        tma_sfa_desc_list.append(
+            _tma.create_tensor_map_tiled_from_view(
+                sfa_fp16_tensor,
+                dtype=cutlass.Uint16,
+                box_dims=(256, sf_tma_box_k, sfa_tma_box_mn, 1),
+                stride_order=(0, 1, 2, 3),
+                swizzle=_tma.TensorMapSwizzle.none,
+            )
+        )
+    tma_sfb_desc_list = []
+    for _sfb_op in _sfb_operands:
+        sfb_fp16_tensor = cute.make_tensor(
+            cute.recast_ptr(_sfb_op.iterator, dtype=cutlass.Float16),
+            cute.make_layout(
+                (256, rest_k, rest_n, num_experts),
+                stride=(
+                    1,
+                    256,
+                    cute.assume(256 * rest_k, 8),
+                    cute.assume(256 * rest_k * rest_n, 8),
+                ),
+            ),
+        )
+        tma_sfb_desc_list.append(
+            _tma.create_tensor_map_tiled_from_view(
+                sfb_fp16_tensor,
+                dtype=cutlass.Uint16,
+                box_dims=(256, sf_tma_box_k, sfb_tma_box_mn, 1),
+                stride_order=(0, 1, 2, 3),
+                swizzle=_tma.TensorMapSwizzle.none,
+            )
+        )
+
+    cluster_m = cluster_shape_mnk[0]
+    cluster_n = cluster_shape_mnk[1]
+    grid_shape = (grid_num_clusters * cluster_m, cluster_n, 1)
+    _kernel(
+        problem_size[0],
+        problem_size[1],
+        problem_size[2],
+        cutlass.Int32(num_experts),
+        cutlass.Int32(num_groups),
+        first_token_offset,
+        a_tma_workspace,
+        # @@INJECT_HOST_KERNEL_DESC_PASS@@
+        # @@INJECT_MOE_HOST_MA_PASS@@
+        # @@INJECT_HOST_TAP_PASS@@
+        # @@INJECT_HOST_REDUCTION_STRIDE_PASS@@
+        # @@INJECT_HOST_AUX_PASS@@
+    ).launch(
+        grid=grid_shape,
+        block=(threads_per_cta, 1, 1),
+        cluster=cluster_shape_mnk,
+        use_pdl=USE_PDL,
+        stream=stream,
+    )
+
+
+@lru_cache(maxsize=None)
+def compile() -> Callable:
+    out_vec_elems = vec_bytes_epi // (cd_dtype.width // 8)
+    ab_stride_elems = 16 // (ab_dtype.width // 8)
+    sym_m = cute.sym_int64()
+    sym_n = cute.sym_int64(divisibility=out_vec_elems)
+    # K tails are supported: the K loop is ceil_div and the TMA descriptor's global K
+    # extent makes a partial box HW zero-filled. The only real K rule is the 16-byte
+    # TMA contiguous-extent one, already gated by _tma_alignment_reject.
+    sym_k = cute.sym_int64()
+    # Packed K extent: same reasoning as sym_k -- no CTA-tile multiple is required.
+    sym_kp = cute.sym_int64()
+    sym_e = cute.sym_int64()
+    sym_g = cute.sym_int64()
+
+    def _make_fake_a():
+        return make_fake_compact_tensor(
+            a_fake_dtype,
+            (sym_m, sym_kp, 1),
+            stride_order=(1, 0, 2),
+            assumed_align=16,
+        )
+
+    def _make_fake_b():
+        return make_fake_compact_tensor(
+            b_fake_dtype,
+            (sym_n, sym_kp, sym_e),
+            stride_order=(0, 1, 2) if b_is_n_major else (1, 0, 2),
+            assumed_align=16,
+        )
+
+    # SF reaches the kernel as a base pointer only; the host rebuilds the
+    # F8_128x4 view from problem_size. Modes 0/1 and all strides carry no
+    # contract; mode 2 keeps its literal plane count (1 for sfa, sym_e for sfb).
+    def _make_fake_sfa():
+        return cute.runtime.make_fake_tensor(
+            sf_cutlass_dtype,
+            (cute.sym_int64(), cute.sym_int64(), 1),
+            stride=(cute.sym_int64(), cute.sym_int64(), cute.sym_int64()),
+            assumed_align=16,
+        )
+
+    def _make_fake_sfb():
+        return cute.runtime.make_fake_tensor(
+            sf_cutlass_dtype,
+            (cute.sym_int64(), cute.sym_int64(), sym_e),
+            stride=(cute.sym_int64(), cute.sym_int64(), cute.sym_int64()),
+            assumed_align=16,
+        )
+
+    fake_first_token_offset = make_fake_compact_tensor(
+        offset_cutlass_dtype,
+        (sym_g,),
+        stride_order=(0,),
+        assumed_align=offset_cutlass_dtype.width // 8,
+    )
+    cluster_m = cluster_shape_mnk[0]
+    cluster_n = cluster_shape_mnk[1]
+    grid_ctas = grid_num_clusters * cluster_m * cluster_n
+    fake_a_tma_workspace = make_fake_compact_tensor(
+        cutlass.Int64,
+        (grid_ctas * num_a_operands * 16,),
+        stride_order=(0,),
+        assumed_align=128,
+    )
+
+    def _sym_operand_strides(is_mn_major: bool) -> tuple:
+        # Operand is permuted to (M|N, K, L): the unit stride is mode 0 when MN-major, mode 1 when K-major, and never reaches TMA.
+        unit = 0 if is_mn_major else 1
+        return tuple(cute.sym_int64() if i == unit else cute.sym_int64(divisibility=ab_stride_elems) for i in range(3))
+
+    sym_a_strides = []
+    for _ in range(num_a_operands):
+        sym_a_strides.extend(_sym_operand_strides(a_is_m_major))
+    sym_b_strides = []
+    for _ in range(num_b_operands):
+        sym_b_strides.extend(_sym_operand_strides(b_is_n_major))
+    # @@INJECT_COMPILE_REDUCTION_STRIDE_DECLS@@
+    # @@INJECT_COMPILE_AB_FAKES@@
+    # @@INJECT_COMPILE_TAP_FAKES@@
+    problem_size = (
+        sym_m,
+        sym_n,
+        sym_k,
+        sym_e,
+        sym_g,
+        *sym_a_strides,
+        *sym_b_strides,
+        # @@INJECT_COMPILE_REDUCTION_STRIDE_SYMBOLS@@
+    )
+    # @@INJECT_COMPILE_AUX_FAKES@@
+    _fake_stream = make_fake_stream(use_tvm_ffi_env_stream=False)
+    return cute.compile(
+        _host,
+        problem_size,
+        fake_first_token_offset,
+        fake_a_tma_workspace,
+        # @@INJECT_COMPILE_AB_PASS@@
+        # @@INJECT_COMPILE_TAP_PASS@@
+        # @@INJECT_COMPILE_AUX_PASS@@
+        stream=_fake_stream,
+        options=frost_compile_options,
+    )
diff --git a/python/cudnn/gemm/frost/kernel_templates/sm107_moe_grouped_block_scale_matmul_fwd_2ctamma.py b/python/cudnn/gemm/frost/kernel_templates/sm107_moe_grouped_block_scale_matmul_fwd_2ctamma.py
new file mode 100644
index 000000000..a7e34b247
--- /dev/null
+++ b/python/cudnn/gemm/frost/kernel_templates/sm107_moe_grouped_block_scale_matmul_fwd_2ctamma.py
@@ -0,0 +1,1442 @@
+# Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+# SPDX-License-Identifier: MIT
+
+"""sm107 cta_group=2 MoE grouped block-scale matmul fwd: grouped persistent
+scheduler + per-group A TMA descriptor patch + block-scaled MMA.
+
+Per routed group g: ``out[fto[g]:fto[g+1]] = deq(token[range]) @ deq(weight[g%E]).T``,
+token/weight FP4/FP8 dequantized by per-block scale factors inside the MMA
+(supports nvfp4 / mxfp4 / mxfp8). 2-CTA MMA cluster pair (cluster2x1 reference
+design): the leader CTA issues the MMA, the follower consumes; both load their
+operand slice. The TMA warp patches A's descriptor on each routed-group change.
+
+The pipeline is the sm100 one; SM 10.7's block-scale MMA reads a **64-byte K**
+per instruction instead of 32, which shows up in exactly two places (both
+driven by injected constants, so the rest of the file stays in lockstep with
+``sm100_moe_grouped_block_scale_matmul_fwd_2ctamma.py``):
+
+  * half as many MMAs per K-tile, each consuming ``sf_scales_per_inst`` scales
+    (8 at K-block 16, 4 at 32 — it follows the BLOCK SIZE, not the scale
+    dtype). When that exceeds the 4 scales one 128x4 utccp atom holds, a scale
+    *word* spans ``word_atoms`` atoms, and the two SF regions then lay them
+    out DIFFERENTLY: SFB atom-major across its N-blocks, SFA block-major.
+    At K-block 32 ``word_atoms == 1`` (identical to sm100).
+  * fp4 rides the OMMA instruction descriptor (``Tcgen05MxOmmaInstrDesc``,
+    K-mode 2 = 128 fp4 elements); mxfp8 stays on ``Tcgen05MxInstrDesc``
+    (K-mode 1 = 64 fp8 elements). Both take the real operand dtype.
+
+The K-tile itself is unchanged (128 bytes), so the grouped scheduler, the
+per-group A-tensormap patch and the per-group-128-padded SF blob layout are
+byte-for-byte the sm100 ones.
+
+Warp layout (8 warps × 32 = 256 threads/CTA):
+  warps 0–3 : epilogue (warp 0 also allocates TMEM)  — setmaxnreg.inc 216
+  warp  4   : MMA driver (leader CTA runs MMA; follower CTA consumes only)  — setmaxnreg.dec 40
+  warp  5   : TMA producer (both CTAs load their slice; per-group A descriptor patch)  — setmaxnreg.dec 40
+  warp  6   : grouped persistent scheduler  — setmaxnreg.dec 40
+  warp  7   : unused donor — setmaxnreg.dec 40, idle to dealloc barrier
+"""
+
+from __future__ import annotations
+
+from functools import lru_cache
+from typing import Callable
+
+import cutlass.experimental.primitives as nvvm
+from cudnn.gemm.frost.kernel_templates._tile_helpers import (
+    copy_tensormap_to_workspace as _copy_tensormap_to_workspace,
+    fence_tensormap_acquire as _fence_tensormap_acquire,
+    fence_tensormap_release as _fence_tensormap_release,
+    moe_swizzle_tile as _moe_swizzle_tile,
+    replace_tensormap_global_address as _replace_tensormap_global_address,
+    replace_tensormap_global_dim_1 as _replace_tensormap_global_dim_1,
+    TENSOR_MAP_QWORDS,
+)
+import cutlass.experimental.cuda.tensor_map as _tma
+import cutlass
+import cutlass.cute as cute
+from cutlass.cute.runtime import make_fake_compact_tensor
+from cutlass.cute.runtime import make_fake_stream
+from cuda.bindings import driver as _cuda
+
+# @@INJECT_TILE_CONSTANTS@@
+
+
+SCHED_STAGES = 2
+SCHED_SLOT_WORDS = 8
+
+USE_PDL = True
+EPI_SMEM_STAGES = 2
+EPI_SYNC_BAR_ID = 1
+TMEM_ALLOC_BARRIER_ID = 2
+
+
+@cute.jit
+def _moe_auto_swizzle_w(group_rows, n, k, nt_n):
+    """N-super-block width for one routed group, resolved per group.
+
+    Same rule as the dense path, but the "M side" is THIS group's token slice, not the
+    whole token tensor: block along the shorter of (group tokens, expert weight), capped
+    by what L2 can hold onto. A group spanning one m-tile makes both orders identical.
+    """
+    if cutlass.const_expr(tile_swizzle_n > 0):
+        return tile_swizzle_n
+    budget = cutlass.Int64(swizzle_l2_budget_bytes)
+    row_bytes = (cutlass.Int64(ab_dtype.width) * k) // 8
+    cap = cutlass.max(budget // (row_bytes * cgrp_tile_mnk[1]), cutlass.Int64(1))
+    w = cutlass.min(cutlass.Int64(nt_n), cap)
+    rows = cutlass.Int64(group_rows)
+    if cutlass.min(rows, n) * row_bytes <= budget and rows <= n:
+        w = cutlass.Int64(1)
+    return cutlass.Int32(w)
+
+
+def _b_collector_op(mi):
+    """B is identical across the M sub-blocks (only A's address advances), so the
+    first MMA fills the B collector and the rest read it back instead of
+    re-fetching the same operand from SMEM."""
+    if cutlass.const_expr(not b_collector_ok or num_mma_m == 1):
+        return None
+    if cutlass.const_expr(mi == 0):
+        return nvvm.Tcgen05MMACollectorOp.FILL
+    if cutlass.const_expr(mi == num_mma_m - 1):
+        return nvvm.Tcgen05MMACollectorOp.LASTUSE
+    return nvvm.Tcgen05MMACollectorOp.USE
+
+
+@cute.kernel
+def _kernel(
+    m: cutlass.Int64,
+    n: cutlass.Int64,
+    k: cutlass.Int64,
+    num_experts: cutlass.Int32,
+    num_groups: cutlass.Int32,
+    first_token_offset: cute.Tensor,
+    a_tma_workspace: cute.Tensor,
+    # @@INJECT_KERNEL_AB_DESC_PARAMS@@
+    # @@INJECT_MOE_KERNEL_MA_PARAMS@@
+    # @@INJECT_KERNEL_TAP_PARAMS@@
+    # @@INJECT_KERNEL_REDUCTION_STRIDE_PARAMS@@
+    # @@INJECT_KERNEL_AUX_PARAMS@@
+) -> None:
+    # @@INJECT_AB_DESC_LISTS@@
+
+    # @@INJECT_MOE_MA_LIST@@
+
+    mma_warp_id = 4
+    tma_warp_id = 5
+    scheduler_warp_id = 6
+    unused_warp_id = 7
+    num_epilogue_warps = 4
+    epi_reg_count = 232
+    prod_reg_count = 24
+
+    warp_idx = cute.arch.warp_idx()
+    warp_idx = cute.arch.make_warp_uniform(warp_idx)
+    elect_one = nvvm.elect_sync()
+
+    tidx = cute.arch.thread_idx()[0]
+    bidx = cute.arch.block_idx()[0]
+    bidy = cute.arch.block_idx()[1]
+    bidz = cute.arch.block_idx()[2]
+    gridx = cute.arch.grid_dim()[0]
+    gridy = cute.arch.grid_dim()[1]
+
+    cluster_m = cluster_shape_mnk[0]
+    cluster_n = cluster_shape_mnk[1]
+    cluster_size = cluster_m * cluster_n * cluster_shape_mnk[2]
+
+    cta_rank_in_cluster = cute.arch.block_idx_in_cluster()
+    m_rank = cta_rank_in_cluster % cluster_m
+    n_rank = cta_rank_in_cluster // cluster_m
+    pair_member = m_rank % 2
+    pair_m_idx = m_rank // 2
+    is_pair_leader = pair_member == 0
+    pair_leader_rank = pair_m_idx * 2 + n_rank * cluster_m
+
+    cluster_linear_init = bidx // cluster_m
+
+    if warp_idx == mma_warp_id:
+        for _i in cutlass.range_constexpr(num_a_operands):
+            nvvm.prefetch_tensormap(tma_a_descs[_i].get_ptr())
+            nvvm.prefetch_tensormap(tma_sfa_descs[_i].get_ptr())
+        for _j in cutlass.range_constexpr(num_b_operands):
+            nvvm.prefetch_tensormap(tma_b_descs[_j].get_ptr())
+            nvvm.prefetch_tensormap(tma_sfb_descs[_j].get_ptr())
+
+    a_pattern = 0
+    for n_idx in cutlass.range_constexpr(cluster_n):
+        a_pattern = a_pattern | (1 << (n_idx * cluster_m))
+    b_pattern = 0
+    for pm_idx in cutlass.range_constexpr(cluster_m // 2):
+        b_pattern = b_pattern | (1 << (pm_idx * 2))
+
+    if cutlass.const_expr(multicast_a):
+        tma_mcast_mask_a = cutlass.Int16(a_pattern << m_rank)
+    else:
+        tma_mcast_mask_a = cutlass.Int16(1 << cta_rank_in_cluster)
+    if cutlass.const_expr(multicast_b):
+        tma_mcast_mask_b = cutlass.Int16((b_pattern << pair_member) << (n_rank * cluster_m))
+    else:
+        tma_mcast_mask_b = cutlass.Int16(1 << cta_rank_in_cluster)
+
+    _smem_sys_reserved = cutlass.Array(cutlass.Int8, 1024, space=cutlass.AddressSpace.smem, alignment=1)
+
+    ab_full_mbar_ptr = cutlass.Array(cutlass.Int64, ab_stages, space=cutlass.AddressSpace.smem)
+    ab_empty_mbar_ptr = cutlass.Array(cutlass.Int64, ab_stages, space=cutlass.AddressSpace.smem)
+    acc_empty_mbar_ptr = cutlass.Array(cutlass.Int64, acc_stages, space=cutlass.AddressSpace.smem)
+    acc_full_mbar_ptr = cutlass.Array(cutlass.Int64, acc_stages, space=cutlass.AddressSpace.smem)
+    tmem_dealloc_mbar_ptr = cutlass.Array(cutlass.Int64, 1, space=cutlass.AddressSpace.smem)
+    tmem_ptr_i32 = cutlass.Array(cutlass.Int32, 1, space=cutlass.AddressSpace.smem)
+
+    sched_storage = cutlass.Array(
+        cutlass.Int32,
+        SCHED_STAGES * SCHED_SLOT_WORDS,
+        space=cutlass.AddressSpace.smem,
+        alignment=16,
+    )
+    sched_full_mbar_ptr = cutlass.Array(cutlass.Int64, SCHED_STAGES, space=cutlass.AddressSpace.smem, alignment=8)
+    sched_empty_mbar_ptr = cutlass.Array(cutlass.Int64, SCHED_STAGES, space=cutlass.AddressSpace.smem, alignment=8)
+    tma_a_desc_smem_list = [
+        cutlass.Array(
+            cutlass.Int64,
+            TENSOR_MAP_QWORDS,
+            space=cutlass.AddressSpace.smem,
+            alignment=128,
+        )
+        for _ in range(num_a_operands)
+    ]
+
+    sA_elems = sA_packed_elems
+    sB_elems = sB_packed_elems
+    smem_a_list = [
+        cutlass.Array(
+            ab_dtype,
+            sA_elems * ab_stages,
+            space=cutlass.AddressSpace.smem,
+            alignment=1024,
+        )
+        for _ in range(num_a_operands)
+    ]
+    smem_b_list = [
+        cutlass.Array(
+            ab_dtype,
+            sB_elems * ab_stages,
+            space=cutlass.AddressSpace.smem,
+            alignment=1024,
+        )
+        for _ in range(num_b_operands)
+    ]
+    smem_sfa_list = [
+        cutlass.Array(
+            cutlass.Uint8,
+            sfa_smem_bytes * ab_stages,
+            space=cutlass.AddressSpace.smem,
+            alignment=1024,
+        )
+        for _ in range(num_a_operands)
+    ]
+    smem_sfb_list = [
+        cutlass.Array(
+            cutlass.Uint8,
+            sfb_smem_bytes * ab_stages,
+            space=cutlass.AddressSpace.smem,
+            alignment=1024,
+        )
+        for _ in range(num_b_operands)
+    ]
+
+    acc_empty_count = num_epilogue_warps * 2
+    cta_group = 2
+    ab_empty_count = (cluster_m // cta_group) + cluster_n - 1
+    sched_empty_count = 1 + 1 + num_epilogue_warps
+    if warp_idx == 0:
+        if cutlass.const_expr(use_acc_overlap):
+            if elect_one:
+                nvvm.mbarrier_init(tmem_dealloc_mbar_ptr, num_epilogue_warps)
+        else:
+            if elect_one:
+                nvvm.mbarrier_init(tmem_dealloc_mbar_ptr, 32)
+        for i in range(ab_stages):
+            if elect_one:
+                nvvm.mbarrier_init(ab_full_mbar_ptr.subview(i), 1)
+            if elect_one:
+                nvvm.mbarrier_init(ab_empty_mbar_ptr.subview(i), ab_empty_count)
+        for i in range(acc_stages):
+            if elect_one:
+                nvvm.mbarrier_init(acc_full_mbar_ptr.subview(i), 1)
+            if elect_one:
+                nvvm.mbarrier_init(acc_empty_mbar_ptr.subview(i), acc_empty_count)
+        for i in range(SCHED_STAGES):
+            if elect_one:
+                nvvm.mbarrier_init(sched_full_mbar_ptr.subview(i), 1)
+            if elect_one:
+                nvvm.mbarrier_init(sched_empty_mbar_ptr.subview(i), sched_empty_count)
+    nvvm.fence_mbarrier_init()
+    nvvm.barrier_cluster_arrive_relaxed()
+
+    sA_bytes = sA_elems * (ab_dtype.width // 8)
+    sB_bytes = sB_elems * (ab_dtype.width // 8)
+    num_tma_copy_bytes = (num_a_operands * (sA_bytes + sfa_smem_bytes) + num_b_operands * (sB_bytes + sfb_smem_bytes)) * 2
+
+    pair_n_size = cgrp_tile_mnk[1] // cluster_n
+    # Per-CTA output rows one MMA-M block covers. The pair splits M, so this is
+    # the per-CTA mma_inst_m — half the instruction's hardware M.
+    epi_rows_per_mma_m = cta_tile_mnk[0] // num_mma_m
+    if cutlass.const_expr(epi_rows_per_mma_m == 64):
+        # cluster-MMA m=128: the pair also splits N, so each CTA drains N/2.
+        cols_per_acc_stage = pair_n_size // 2
+    else:
+        cols_per_acc_stage = pair_n_size
+    tmem_alloc_bar_count = (num_epilogue_warps + 1) * 32
+
+    nvvm.barrier_cluster_wait()
+    nvvm.barrier_cta_sync(0)
+
+    # @@INJECT_TAP_PTRS@@
+
+    VEC_BYTES = vec_bytes_epi
+    vsize = (VEC_BYTES * 8) // cd_dtype.width
+
+    M = m
+    N = n
+    clusters_along_n = cute.ceil_div(cutlass.Int32(N), cgrp_tile_mnk[1])
+    num_k_tiles = cute.ceil_div(k, cgrp_tile_mnk[2])
+    first_token_arr = cutlass.make_array_view(first_token_offset)
+
+    if warp_idx == scheduler_warp_id:
+        nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
+        full_warp_mask = 0xFFFFFFFF
+        shfl_idx_clamp = 0x1F
+        shfl_up_clamp = 0
+        lane = cute.arch.lane_idx()
+        gemm_s = cutlass.Int32(M)
+        sched_stage = cutlass.Int32(0)
+        sched_empty_phase = cutlass.Int32(1)
+        linear_idx = cutlass.Int32(cluster_linear_init)
+        start_linear_idx = cutlass.Int32(0)
+        total_tiles = cutlass.Int32(0)
+        start_sf_block_m = cutlass.Int32(0)
+        total_sf_blocks_m = cutlass.Int32(0)
+        group_idx = cutlass.Int32(0)
+        is_tile_valid = cutlass.Int32(1)
+        cached_next_end = cutlass.Int32(0)
+        if lane + 1 < num_groups:
+            cached_next_end = cutlass.Int32(first_token_arr[lane + 1])
+        else:
+            cached_next_end = gemm_s
+        tile_lower_bound = nvvm.shfl_sync(full_warp_mask, cached_next_end, 1, shfl_up_clamp, nvvm.Shfl.UP)
+        cached_next_begin = cutlass.Int32(0)
+        if lane != 0:
+            cached_next_begin = tile_lower_bound
+
+        while is_tile_valid != 0:
+            group_begin = cached_next_begin
+            group_end = cached_next_end
+
+            if linear_idx >= start_linear_idx + total_tiles:
+                group_idx += lane
+                is_search_live = cutlass.Int32(1)
+                while is_search_live != 0:
+                    cached_group_begin = cached_next_begin
+                    cached_group_end = cached_next_end
+                    tile_start_idx = nvvm.shfl_sync(
+                        full_warp_mask,
+                        cached_next_end,
+                        31,
+                        shfl_idx_clamp,
+                        nvvm.Shfl.IDX,
+                    )
+                    next_end_group = group_idx + 32 + 1
+                    if next_end_group < num_groups:
+                        cached_next_end = cutlass.Int32(first_token_arr[next_end_group])
+                    else:
+                        cached_next_end = gemm_s
+                    tile_lower_bound = nvvm.shfl_sync(
+                        full_warp_mask,
+                        cached_next_end,
+                        1,
+                        shfl_up_clamp,
+                        nvvm.Shfl.UP,
+                    )
+                    if lane != 0:
+                        cached_next_begin = tile_lower_bound
+                    else:
+                        cached_next_begin = tile_start_idx
+
+                    group_m = cached_group_end - cached_group_begin
+                    total_tiles = cute.ceil_div(group_m, cgrp_tile_mnk[0]) * clusters_along_n
+                    total_sf_blocks_m = cute.ceil_div(group_m, 128)
+                    prefix_tiles = total_tiles
+                    prefix_sf = total_sf_blocks_m
+                    for delta in (1, 2, 4, 8, 16):
+                        prefix_delta = nvvm.shfl_sync(
+                            full_warp_mask,
+                            prefix_tiles,
+                            delta,
+                            shfl_up_clamp,
+                            nvvm.Shfl.UP,
+                        )
+                        prefix_sf_delta = nvvm.shfl_sync(
+                            full_warp_mask,
+                            prefix_sf,
+                            delta,
+                            shfl_up_clamp,
+                            nvvm.Shfl.UP,
+                        )
+                        if lane >= delta:
+                            prefix_tiles += prefix_delta
+                            prefix_sf += prefix_sf_delta
+                    start_linear_idx += prefix_tiles - total_tiles
+                    start_sf_block_m += prefix_sf - total_sf_blocks_m
+                    thread_succeed = nvvm.vote_sync(
+                        full_warp_mask,
+                        linear_idx < start_linear_idx + total_tiles,
+                        nvvm.VoteSync.BALLOT,
+                    )
+                    if thread_succeed != 0:
+                        winning_lane = cutlass.Int32(31) - cute.arch.bfind(cute.arch.brev(thread_succeed)).to(cutlass.Int32)
+                        group_idx = nvvm.shfl_sync(
+                            full_warp_mask,
+                            group_idx,
+                            winning_lane,
+                            shfl_idx_clamp,
+                            nvvm.Shfl.IDX,
+                        )
+                        start_linear_idx = nvvm.shfl_sync(
+                            full_warp_mask,
+                            start_linear_idx,
+                            winning_lane,
+                            shfl_idx_clamp,
+                            nvvm.Shfl.IDX,
+                        )
+                        total_tiles = nvvm.shfl_sync(
+                            full_warp_mask,
+                            total_tiles,
+                            winning_lane,
+                            shfl_idx_clamp,
+                            nvvm.Shfl.IDX,
+                        )
+                        start_sf_block_m = nvvm.shfl_sync(
+                            full_warp_mask,
+                            start_sf_block_m,
+                            winning_lane,
+                            shfl_idx_clamp,
+                            nvvm.Shfl.IDX,
+                        )
+                        tile_start_idx = nvvm.shfl_sync(
+                            full_warp_mask,
+                            cached_group_begin,
+                            winning_lane,
+                            shfl_idx_clamp,
+                            nvvm.Shfl.IDX,
+                        )
+                        group_end_idx = group_idx + lane + 1
+                        if group_end_idx < num_groups:
+                            cached_next_end = cutlass.Int32(first_token_arr[group_end_idx])
+                        else:
+                            cached_next_end = gemm_s
+                        tile_lower_bound = nvvm.shfl_sync(
+                            full_warp_mask,
+                            cached_next_end,
+                            1,
+                            shfl_up_clamp,
+                            nvvm.Shfl.UP,
+                        )
+                        if lane != 0:
+                            cached_next_begin = tile_lower_bound
+                        else:
+                            cached_next_begin = tile_start_idx
+                        group_begin = cached_next_begin
+                        group_end = cached_next_end
+                        is_search_live = cutlass.Int32(0)
+                    else:
+                        group_idx += 32
+                        first_lane_group = nvvm.shfl_sync(
+                            full_warp_mask,
+                            group_idx,
+                            0,
+                            shfl_idx_clamp,
+                            nvvm.Shfl.IDX,
+                        )
+                        if first_lane_group >= num_groups:
+                            is_tile_valid = cutlass.Int32(0)
+                            is_search_live = cutlass.Int32(0)
+                        else:
+                            next_start_linear_idx = start_linear_idx + total_tiles
+                            start_linear_idx = nvvm.shfl_sync(
+                                full_warp_mask,
+                                next_start_linear_idx,
+                                31,
+                                shfl_idx_clamp,
+                                nvvm.Shfl.IDX,
+                            )
+                            next_start_sf = start_sf_block_m + total_sf_blocks_m
+                            start_sf_block_m = nvvm.shfl_sync(
+                                full_warp_mask,
+                                next_start_sf,
+                                31,
+                                shfl_idx_clamp,
+                                nvvm.Shfl.IDX,
+                            )
+
+            coord_expert = cutlass.Int32(0)
+            cluster_tile_m = cutlass.Int32(0)
+            coord_n = cutlass.Int32(0)
+            if is_tile_valid != 0:
+                local_linear_idx = linear_idx - start_linear_idx
+                group_nt_m = total_tiles // clusters_along_n
+                cluster_tile_m, coord_n = _moe_swizzle_tile(
+                    local_linear_idx,
+                    group_nt_m,
+                    clusters_along_n,
+                    _moe_auto_swizzle_w(group_nt_m * cgrp_tile_mnk[0], N, k, clusters_along_n),
+                )
+                coord_expert = group_idx % num_experts
+                linear_idx += grid_num_clusters
+
+            while not nvvm.mbarrier_try_wait_parity(
+                sched_empty_mbar_ptr.subview(sched_stage),
+                sched_empty_phase,
+                time_limit=10_000_000,
+            ):
+                pass
+            if lane == 0:
+                slot = sched_storage.subview(sched_stage * SCHED_SLOT_WORDS)
+                (slot.subview(0)).store(coord_expert)
+                (slot.subview(1)).store(cluster_tile_m)
+                (slot.subview(2)).store(coord_n)
+                (slot.subview(3)).store(is_tile_valid)
+                (slot.subview(4)).store(group_begin)
+                (slot.subview(5)).store(group_end)
+                (slot.subview(6)).store(start_sf_block_m)
+                (slot.subview(7)).store(group_idx)
+                nvvm.mbarrier_arrive(sched_full_mbar_ptr.subview(sched_stage))
+
+            sched_stage += 1
+            if sched_stage == SCHED_STAGES:
+                sched_stage = cutlass.Int32(0)
+                sched_empty_phase = sched_empty_phase ^ 1
+
+    if warp_idx == tma_warp_id:
+        nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
+        if cutlass.const_expr(USE_PDL):
+            if elect_one:
+                nvvm.griddepcontrol("wait")
+        ab_empty_phase_bit = cutlass.Int32(1)
+        ab_iter = cutlass.Int32(0)
+        sched_stage = cutlass.Int32(0)
+        sched_full_phase = cutlass.Int32(0)
+        is_valid = cutlass.Int32(1)
+        logical_cta_tile_n = cgrp_tile_mnk[1] // cluster_n
+
+        lane = tidx % 32
+        block_linear = bidx + bidy * gridx
+        cta_desc_base_list = [a_tma_workspace.iterator.raw_ptr() + (block_linear * num_a_operands + _ai) * TENSOR_MAP_QWORDS for _ai in range(num_a_operands)]
+        a_desc_tma_ptr_list = [
+            cute.make_ptr(
+                cutlass.Int64,
+                cta_desc_base_list[_ai].toint(),
+                mem_space=cute.AddressSpace.generic,
+            )
+            for _ai in range(num_a_operands)
+        ]
+        previous_group_begin = cutlass.Int32(-1)
+        if elect_one:
+            for _ai in cutlass.range_constexpr(num_a_operands):
+                _copy_tensormap_to_workspace(tma_a_descs[_ai].get_ptr(), tma_a_desc_smem_list[_ai])
+        nvvm.bar_warp_sync(0xFFFFFFFF)
+
+        while is_valid != 0:
+            while not nvvm.mbarrier_try_wait_parity(
+                sched_full_mbar_ptr.subview(sched_stage),
+                sched_full_phase,
+                time_limit=10_000_000,
+            ):
+                pass
+            slot = sched_storage.subview(sched_stage * SCHED_SLOT_WORDS)
+            coord_expert = (slot.subview(0)).load()
+            tile_m = (slot.subview(1)).load()
+            tile_n = (slot.subview(2)).load()
+            is_valid = (slot.subview(3)).load()
+            group_begin = (slot.subview(4)).load()
+            group_end = (slot.subview(5)).load()
+            start_sf_block_m = (slot.subview(6)).load()
+            if elect_one:
+                nvvm.mbarrier_arrive(sched_empty_mbar_ptr.subview(sched_stage))
+            sched_stage += 1
+            if sched_stage == SCHED_STAGES:
+                sched_stage = cutlass.Int32(0)
+                sched_full_phase = sched_full_phase ^ 1
+
+            if is_valid != 0:
+                coord_m_group = tile_m * cgrp_tile_mnk[0] + m_rank * cta_tile_mnk[0]
+                coord_n_per_cta = tile_n * cgrp_tile_mnk[1] + n_rank * logical_cta_tile_n + pair_member * cta_tile_mnk[1]
+                coord_n_pair = tile_n * cgrp_tile_mnk[1] + n_rank * logical_cta_tile_n
+                sfa_m_block = start_sf_block_m + coord_m_group // 128
+                sfb_n_block = coord_n_pair // 128
+
+                if group_begin != previous_group_begin:
+                    previous_group_begin = group_begin
+                    for _ai in cutlass.range_constexpr(num_a_operands):
+                        _fence_tensormap_acquire(a_desc_tma_ptr_list[_ai])
+                    for _ai in cutlass.range_constexpr(num_a_operands):
+                        if elect_one:
+                            row_base = mA_list[_ai].iterator.raw_ptr().toint() + ((group_begin * a_stride_m_list[_ai] * ab_dtype.width) >> 3)
+                            _replace_tensormap_global_address(tma_a_desc_smem_list[_ai], row_base)
+                            _replace_tensormap_global_dim_1(tma_a_desc_smem_list[_ai], group_end - group_begin)
+                        nvvm.bar_warp_sync(0xFFFFFFFF)
+                        if lane < TENSOR_MAP_QWORDS:
+                            (cta_desc_base_list[_ai] + lane).store((tma_a_desc_smem_list[_ai].subview(lane)).load())
+                        nvvm.bar_warp_sync(0xFFFFFFFF)
+                        _fence_tensormap_release()
+
+                for k_tile_idx in range(num_k_tiles):
+                    stage = ab_iter % ab_stages
+                    if stage == 0 and ab_iter != 0:
+                        ab_empty_phase_bit = ab_empty_phase_bit ^ 1
+
+                    while not nvvm.mbarrier_try_wait_parity(
+                        ab_empty_mbar_ptr.subview(stage),
+                        ab_empty_phase_bit,
+                        time_limit=10_000_000,
+                    ):
+                        pass
+
+                    coord_k = k_tile_idx * cgrp_tile_mnk[2]
+                    coord_sf_k = k_tile_idx * sf_tma_box_k
+
+                    if is_pair_leader:
+                        if elect_one:
+                            nvvm.mbarrier_arrive_expect_tx(ab_full_mbar_ptr.subview(stage), num_tma_copy_bytes)
+                    a_issue = (not multicast_a) or (n_rank == 0)
+                    b_issue = (not multicast_b) or (pair_m_idx == 0)
+                    if a_issue:
+                        for _ai in cutlass.range_constexpr(num_a_operands):
+                            if elect_one:
+                                nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                    smem_a_list[_ai].subview(sA_elems * stage),
+                                    a_desc_tma_ptr_list[_ai],
+                                    (coord_k, coord_m_group, cutlass.Int32(0)),
+                                    ab_full_mbar_ptr.subview(stage),
+                                    [],
+                                    multicast_mask=tma_mcast_mask_a,
+                                    group=nvvm.CTAGroup.CTA_2,
+                                )
+                        for _ai in cutlass.range_constexpr(num_a_operands):
+                            if elect_one:
+                                nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                    smem_sfa_list[_ai].subview(sfa_smem_bytes * stage),
+                                    tma_sfa_descs[_ai].get_ptr(),
+                                    (0, coord_sf_k, sfa_m_block, cutlass.Int32(0)),
+                                    ab_full_mbar_ptr.subview(stage),
+                                    [],
+                                    multicast_mask=tma_mcast_mask_a,
+                                    group=nvvm.CTAGroup.CTA_2,
+                                )
+                    if b_issue:
+                        for _bj in cutlass.range_constexpr(num_b_operands):
+                            sB_stage = smem_b_list[_bj].subview(sB_elems * stage)
+                            if cutlass.const_expr(b_is_n_major):
+                                for n_group in cutlass.range_constexpr(cta_tile_mnk[1] // b_tma_group_elems):
+                                    if elect_one:
+                                        nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                            sB_stage.subview(n_group * b_tma_group_elems * cgrp_tile_mnk[2]),
+                                            tma_b_descs[_bj].get_ptr(),
+                                            (
+                                                coord_n_per_cta + n_group * b_tma_group_elems,
+                                                coord_k,
+                                                coord_expert,
+                                            ),
+                                            ab_full_mbar_ptr.subview(stage),
+                                            [],
+                                            multicast_mask=tma_mcast_mask_b,
+                                            group=nvvm.CTAGroup.CTA_2,
+                                        )
+                            else:
+                                if elect_one:
+                                    nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                        sB_stage,
+                                        tma_b_descs[_bj].get_ptr(),
+                                        (coord_k, coord_n_per_cta, coord_expert),
+                                        ab_full_mbar_ptr.subview(stage),
+                                        [],
+                                        multicast_mask=tma_mcast_mask_b,
+                                        group=nvvm.CTAGroup.CTA_2,
+                                    )
+                        for _bj in cutlass.range_constexpr(num_b_operands):
+                            if elect_one:
+                                nvvm.cp_async_bulk_tensor_shared_cluster_global(
+                                    smem_sfb_list[_bj].subview(sfb_smem_bytes * stage),
+                                    tma_sfb_descs[_bj].get_ptr(),
+                                    (0, coord_sf_k, sfb_n_block, coord_expert),
+                                    ab_full_mbar_ptr.subview(stage),
+                                    [],
+                                    multicast_mask=tma_mcast_mask_b,
+                                    group=nvvm.CTAGroup.CTA_2,
+                                )
+
+                    ab_iter += 1
+
+        tail_stage = ab_iter % ab_stages
+        tail_phase = ab_empty_phase_bit
+        if tail_stage == 0 and ab_iter != 0:
+            tail_phase = tail_phase ^ 1
+        for _ in range(ab_stages - 1):
+            tail_stage = tail_stage + 1
+            if tail_stage == ab_stages:
+                tail_stage = cutlass.Int32(0)
+                tail_phase = tail_phase ^ 1
+        if elect_one:
+            while not nvvm.mbarrier_try_wait_parity(ab_empty_mbar_ptr.subview(tail_stage), tail_phase, time_limit=10_000_000):
+                pass
+
+    pair_mask = cutlass.Int16(3) << pair_leader_rank
+    a_arrive_pattern = 0
+    for n_idx in cutlass.range_constexpr(cluster_n):
+        a_arrive_pattern = a_arrive_pattern | (1 << (n_idx * cluster_m))
+    b_arrive_pattern = 0
+    for m_idx in cutlass.range_constexpr(cluster_m):
+        b_arrive_pattern = b_arrive_pattern | (1 << m_idx)
+    a_part = a_arrive_pattern << m_rank
+    a_part = a_part | (a_part << 1)
+    b_part = b_arrive_pattern << (n_rank * cluster_m)
+    ab_empty_arrive_mask = cutlass.Int16(a_part | b_part)
+    if warp_idx == mma_warp_id:
+        nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
+        nvvm.tcgen05_alloc(
+            tmem_ptr_i32,
+            cutlass.Int32(num_tmem_alloc_cols),
+            is_exclusive=tmem_alloc_exclusive,
+            group=nvvm.CTAGroup.CTA_2,
+        )
+        nvvm.bar_warp_sync(0xFFFFFFFF)
+        nvvm.barrier_cta_arrive(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
+        tmem_raw_addr = tmem_ptr_i32.load()
+        base_col_id_root = tmem_raw_addr & 0xFFFF
+        base_row_id = tmem_raw_addr >> 16
+        peer_cta_rank = cta_rank_in_cluster ^ 1
+        if is_pair_leader:
+            ab_full_phase_bit = cutlass.Int32(0)
+            ab_iter = cutlass.Int32(0)
+            acc_empty_phase_bit = cutlass.Int32(1)
+            tile_iter = cutlass.Int32(0)
+            is_valid = cutlass.Int32(1)
+            sched_stage = cutlass.Int32(0)
+            sched_full_phase = cutlass.Int32(0)
+            acc_stage = cutlass.Int32(0)
+            # fp4 packs its K-mode into the OMMA descriptor's 2-bit split field; fp8
+            # keeps the MX descriptor's 1-bit one. Both are built once, outside the
+            # loops — the fields depend only on j (the scale id within a word).
+            if cutlass.const_expr(idesc_is_omma):
+                idesc_by_j = [
+                    cutlass.experimental.primitives.Tcgen05MxOmmaInstrDesc.build(
+                        a_dtype=idesc_a_dtype,
+                        b_dtype=idesc_b_dtype,
+                        scale_format=sf_scale_format,
+                        n_dim=mma_n_dim,
+                        m_dim=mma_m_dim,
+                        a_major=mma_a_major,
+                        b_major=mma_b_major,
+                        a_sf_id=j * sf_scales_per_inst,
+                        b_sf_id=j * sf_scales_per_inst,
+                        k_dim=mma_k_dim_mode,
+                    )
+                    for j in range(sf_insts_per_atom)
+                ]
+            else:
+                idesc_by_j = [
+                    cutlass.experimental.primitives.Tcgen05MxInstrDesc.build(
+                        a_dtype=idesc_a_dtype,
+                        b_dtype=idesc_b_dtype,
+                        scale_format=sf_scale_format,
+                        n_dim=mma_n_dim,
+                        m_dim=mma_m_dim,
+                        a_major=mma_a_major,
+                        b_major=mma_b_major,
+                        a_sf_id=j * sf_scales_per_inst,
+                        b_sf_id=j * sf_scales_per_inst,
+                        k_dim=mma_k_dim_mode,
+                    )
+                    for j in range(sf_insts_per_atom)
+                ]
+            sfa_tmem_bases = [(base_row_id << 16) | (base_col_id_root + sfa_col_bases[i]) for i in range(num_a_operands)]
+            sfb_tmem_bases = [(base_row_id << 16) | (base_col_id_root + sfb_col_bases[j]) for j in range(num_b_operands)]
+            s2t_shape, s2t_multicast = nvvm.S2TCopyMode.S2T_32x128b_WARPX4
+            sfa_scale_ptrs = [nvvm.make_tmem_ptr(b, cutlass.Float32) for b in sfa_tmem_bases]
+            sfb_scale_ptrs = [nvvm.make_tmem_ptr(b, cutlass.Float32) for b in sfb_tmem_bases]
+            # utccp destination per (MN-block, atom within the scale word). SFB
+            # is atom-MAJOR across the N-blocks because ONE instruction walks
+            # all of them; SFA is block-major because one instruction covers
+            # exactly one 128-row block, so that word has to be contiguous.
+            # Both collapse to the same addresses at a single block, and to
+            # sm100's layout at word_atoms == 1.
+            sfa_dst_ptrs = [
+                [
+                    [nvvm.make_tmem_ptr(sfa_tmem_bases[i] + m * registers_per_block + a * registers_per_atom, cutlass.Float32) for a in range(word_atoms)]
+                    for m in range(num_blocks_m)
+                ]
+                for i in range(num_a_operands)
+            ]
+            sfb_dst_ptrs = [
+                [
+                    [nvvm.make_tmem_ptr(sfb_tmem_bases[j] + (a * num_blocks_n + m) * registers_per_atom, cutlass.Float32) for a in range(word_atoms)]
+                    for m in range(num_blocks_n)
+                ]
+                for j in range(num_b_operands)
+            ]
+            while is_valid != 0:
+                while not nvvm.mbarrier_try_wait_parity(
+                    sched_full_mbar_ptr.subview(sched_stage),
+                    sched_full_phase,
+                    time_limit=10_000_000,
+                ):
+                    pass
+                is_valid = (sched_storage.subview(sched_stage * SCHED_SLOT_WORDS).subview(3)).load()
+                if elect_one:
+                    nvvm.mbarrier_arrive(sched_empty_mbar_ptr.subview(sched_stage))
+                sched_stage += 1
+                if sched_stage == SCHED_STAGES:
+                    sched_stage = cutlass.Int32(0)
+                    sched_full_phase = sched_full_phase ^ 1
+
+                if is_valid != 0:
+                    acc_stage = tile_iter % acc_stages
+                    if acc_stage == 0 and tile_iter != 0:
+                        acc_empty_phase_bit = acc_empty_phase_bit ^ 1
+
+                    while not nvvm.mbarrier_try_wait_parity(
+                        acc_empty_mbar_ptr.subview(acc_stage),
+                        acc_empty_phase_bit,
+                        time_limit=10_000_000,
+                    ):
+                        pass
+
+                    if cutlass.const_expr(use_acc_overlap):
+                        acc_base_col = base_col_id_root + (tile_iter % 2) * acc_stage_stride
+                    else:
+                        acc_base_col = base_col_id_root + acc_stage * acc_region_cols
+                    # One accumulator per (gemm, M block); M block mi sits
+                    # epi_cols_per_mma_m columns further into its GEMM's region and
+                    # reads SF word block mi (SF words are one per 128 rows).
+                    acc_tmem_ptrs = [
+                        [
+                            nvvm.make_tmem_ptr(
+                                (base_row_id << 16) | (acc_base_col + g * acc_gemm_stride + mi * epi_cols_per_mma_m),
+                                cutlass.Float32,
+                            )
+                            for mi in range(num_mma_m)
+                        ]
+                        for g in range(num_gemms)
+                    ]
+
+                    scale_d = cutlass.Boolean(False)
+                    for k_tile_idx in range(num_k_tiles):
+                        stage = ab_iter % ab_stages
+                        if stage == 0 and ab_iter != 0:
+                            ab_full_phase_bit = ab_full_phase_bit ^ 1
+
+                        while not nvvm.mbarrier_try_wait_parity(
+                            ab_full_mbar_ptr.subview(stage),
+                            ab_full_phase_bit,
+                            time_limit=10_000_000,
+                        ):
+                            pass
+
+                        desc_a_bases = [
+                            cutlass.experimental.primitives.Tcgen05SmemDesc.build(
+                                start_address=smem_a_list[i].subview(sA_elems * stage),
+                                leading_byte_offset=a_smem_desc_leading_byte_offset,
+                                stride_byte_offset=a_smem_desc_stride_byte_offset,
+                                layout=ab_smem_swizzle,
+                            )
+                            for i in range(num_a_operands)
+                        ]
+                        desc_b_bases = [
+                            cutlass.experimental.primitives.Tcgen05SmemDesc.build(
+                                start_address=smem_b_list[j].subview(sB_elems * stage),
+                                leading_byte_offset=b_smem_desc_leading_byte_offset,
+                                stride_byte_offset=b_smem_desc_stride_byte_offset,
+                                layout=ab_smem_swizzle,
+                            )
+                            for j in range(num_b_operands)
+                        ]
+                        desc_sfa_bases = [
+                            cutlass.experimental.primitives.Tcgen05SmemDesc.build(
+                                start_address=smem_sfa_list[i].subview(sfa_smem_bytes * stage),
+                                leading_byte_offset=16,
+                                stride_byte_offset=128,
+                                layout=cutlass.experimental.primitives.Tcgen05SmemSwizzle.NONE,
+                            )
+                            for i in range(num_a_operands)
+                        ]
+                        desc_sfb_bases = [
+                            cutlass.experimental.primitives.Tcgen05SmemDesc.build(
+                                start_address=smem_sfb_list[j].subview(sfb_smem_bytes * stage),
+                                leading_byte_offset=16,
+                                stride_byte_offset=128,
+                                layout=cutlass.experimental.primitives.Tcgen05SmemSwizzle.NONE,
+                            )
+                            for j in range(num_b_operands)
+                        ]
+
+                        # One SF word per group of MMAs, refreshed right before they
+                        # read it. A word spans word_atoms consecutive K-atoms in SMEM.
+                        for atom_r in cutlass.range(num_sf_atoms, unroll_full=True):
+                            for _ai in cutlass.range_constexpr(num_a_operands):
+                                for _m in cutlass.range_constexpr(num_blocks_m):
+                                    for _a in cutlass.range_constexpr(word_atoms):
+                                        if elect_one:
+                                            nvvm.tcgen05_cp(
+                                                s2t_shape,
+                                                sfa_dst_ptrs[_ai][_m][_a],
+                                                desc_sfa_bases[_ai] + (sf_atom_desc_stride * (atom_r * word_atoms + _a) + sf_block_desc_stride * _m),
+                                                group=nvvm.CTAGroup.CTA_2,
+                                                multicast=s2t_multicast,
+                                            )
+                            for _bj in cutlass.range_constexpr(num_b_operands):
+                                for _m in cutlass.range_constexpr(num_blocks_n):
+                                    for _a in cutlass.range_constexpr(word_atoms):
+                                        if elect_one:
+                                            nvvm.tcgen05_cp(
+                                                s2t_shape,
+                                                sfb_dst_ptrs[_bj][_m][_a],
+                                                desc_sfb_bases[_bj] + (sf_atom_desc_stride * (atom_r * word_atoms + _a) + sf_block_desc_stride * _m),
+                                                group=nvvm.CTAGroup.CTA_2,
+                                                multicast=s2t_multicast,
+                                            )
+                            for j in cutlass.range_constexpr(sf_insts_per_atom):
+                                k_block_idx = atom_r * sf_insts_per_atom + j
+                                idesc_k = idesc_by_j[j]
+                                for g in cutlass.range_constexpr(num_gemms):
+                                    _ai = gemm_a_idx[g]
+                                    _bj = gemm_b_idx[g]
+                                    desc_a_k = desc_a_bases[_ai].advance_start_address(a_smem_k_step_bytes * k_block_idx)
+                                    desc_b = desc_b_bases[_bj].advance_start_address(b_smem_k_step_bytes * k_block_idx)
+                                    for mi in cutlass.range_constexpr(num_mma_m):
+                                        # The M sub-block offset is a whole SMEM swizzle atom, so
+                                        # the descriptor's swizzle phase is preserved. B and its SF
+                                        # are shared; A's SF word block follows the M block.
+                                        desc_a = desc_a_k.advance_start_address(a_smem_m_step_bytes * mi)
+                                        if elect_one:
+                                            nvvm.tcgen05_mma_block_scale(
+                                                mma_block_scale_kind,
+                                                nvvm.CTAGroup.CTA_2,
+                                                acc_tmem_ptrs[g][mi],
+                                                desc_a,
+                                                desc_b,
+                                                idesc_k,
+                                                enable_input_d=scale_d,
+                                                scale_a=sfa_dst_ptrs[_ai][mi][0],
+                                                scale_b=sfb_scale_ptrs[_bj],
+                                                scale_vec_size=scale_vec_size,
+                                                b_collector_op=_b_collector_op(mi),
+                                            )
+                                # Every accumulator sees scale_d=False on exactly the first
+                                # k_block of the tile, so the flip stays outside mi.
+                                scale_d = cutlass.Boolean(True)
+
+                        if elect_one:
+                            nvvm.tcgen05_commit(
+                                ab_empty_mbar_ptr.subview(stage),
+                                multicast_mask=ab_empty_arrive_mask,
+                                group=nvvm.CTAGroup.CTA_2,
+                            )
+                        ab_iter += 1
+
+                    if elect_one:
+                        nvvm.tcgen05_commit(
+                            acc_full_mbar_ptr.subview(acc_stage),
+                            multicast_mask=pair_mask,
+                            group=nvvm.CTAGroup.CTA_2,
+                        )
+                    tile_iter += 1
+
+            if cutlass.const_expr(USE_PDL):
+                if elect_one:
+                    nvvm.griddepcontrol("launch_dependents")
+
+            if tile_iter != 0:
+                tail_stage = acc_stage
+                tail_phase = acc_empty_phase_bit
+                if elect_one:
+                    for _ in range(acc_stages):
+                        tail_stage = tail_stage + 1
+                        if tail_stage == acc_stages:
+                            tail_stage = cutlass.Int32(0)
+                            tail_phase = tail_phase ^ 1
+                        while not nvvm.mbarrier_try_wait_parity(
+                            acc_empty_mbar_ptr.subview(tail_stage),
+                            tail_phase,
+                            time_limit=10_000_000,
+                        ):
+                            pass
+            nvvm.bar_warp_sync(0xFFFFFFFF)
+
+            nvvm.tcgen05_relinquish_alloc_permit(group=nvvm.CTAGroup.CTA_2)
+            peer_mbar = nvvm.mapa(tmem_dealloc_mbar_ptr, peer_cta_rank)
+            while not nvvm.mbarrier_try_wait_parity(tmem_dealloc_mbar_ptr, 0, time_limit=10_000_000):
+                pass
+            if cutlass.const_expr(not use_acc_overlap):
+                nvvm.mbarrier_arrive(peer_mbar, scope=nvvm.MemScope.CLUSTER, relaxed=True)
+            alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
+            nvvm.tcgen05_dealloc(
+                alloc_ptr,
+                cutlass.Int32(num_tmem_alloc_cols),
+                is_exclusive=tmem_alloc_exclusive,
+                group=nvvm.CTAGroup.CTA_2,
+            )
+        else:
+            is_valid = cutlass.Int32(1)
+            sched_stage = cutlass.Int32(0)
+            sched_full_phase = cutlass.Int32(0)
+            while is_valid != 0:
+                while not nvvm.mbarrier_try_wait_parity(
+                    sched_full_mbar_ptr.subview(sched_stage),
+                    sched_full_phase,
+                    time_limit=10_000_000,
+                ):
+                    pass
+                is_valid = (sched_storage.subview(sched_stage * SCHED_SLOT_WORDS).subview(3)).load()
+                if elect_one:
+                    nvvm.mbarrier_arrive(sched_empty_mbar_ptr.subview(sched_stage))
+                sched_stage += 1
+                if sched_stage == SCHED_STAGES:
+                    sched_stage = cutlass.Int32(0)
+                    sched_full_phase = sched_full_phase ^ 1
+
+            if cutlass.const_expr(USE_PDL):
+                if elect_one:
+                    nvvm.griddepcontrol("launch_dependents")
+
+            nvvm.tcgen05_relinquish_alloc_permit(group=nvvm.CTAGroup.CTA_2)
+            peer_mbar = nvvm.mapa(tmem_dealloc_mbar_ptr, peer_cta_rank)
+            if cutlass.const_expr(not use_acc_overlap):
+                nvvm.mbarrier_arrive(peer_mbar, scope=nvvm.MemScope.CLUSTER, relaxed=True)
+            while not nvvm.mbarrier_try_wait_parity(tmem_dealloc_mbar_ptr, 0, time_limit=10_000_000):
+                pass
+            alloc_ptr = cutlass.inttoptr(tmem_raw_addr, 6, cutlass.Int32)
+            nvvm.tcgen05_dealloc(
+                alloc_ptr,
+                cutlass.Int32(num_tmem_alloc_cols),
+                is_exclusive=tmem_alloc_exclusive,
+                group=nvvm.CTAGroup.CTA_2,
+            )
+
+    if warp_idx < num_epilogue_warps:
+        nvvm.setmaxregister(epi_reg_count, nvvm.SetMaxRegisterAction.INCREASE)
+        nvvm.barrier_cta_sync(barrier_id=TMEM_ALLOC_BARRIER_ID, thread_count=tmem_alloc_bar_count)
+        tmem_raw_addr = tmem_ptr_i32.load()
+        base_col_id_root = tmem_raw_addr & 0xFFFF
+        base_row_id = tmem_raw_addr >> 16
+        if cutlass.const_expr(USE_PDL):
+            nvvm.griddepcontrol("wait")
+        tile_iter = cutlass.Int32(0)
+        acc_full_phase_bit = cutlass.Int32(0)
+        is_valid = cutlass.Int32(1)
+        sched_stage = cutlass.Int32(0)
+        sched_full_phase = cutlass.Int32(0)
+
+        row_id_with_warp_offset = base_row_id + warp_idx * 32
+        if cutlass.const_expr(cols_per_acc_stage >= 32):
+            t2r_inst_repx = 32
+            subtile_cnt = cols_per_acc_stage // 32
+        else:
+            t2r_inst_repx = cols_per_acc_stage
+            subtile_cnt = 1
+        shape = nvvm.Tcgen05LdStShape.SHAPE_32X32B
+        lane = tidx % 32
+
+        while is_valid != 0:
+            while not nvvm.mbarrier_try_wait_parity(
+                sched_full_mbar_ptr.subview(sched_stage),
+                sched_full_phase,
+                time_limit=10_000_000,
+            ):
+                pass
+            slot = sched_storage.subview(sched_stage * SCHED_SLOT_WORDS)
+            tile_m = (slot.subview(1)).load()
+            tile_n = (slot.subview(2)).load()
+            is_valid = (slot.subview(3)).load()
+            group_begin = (slot.subview(4)).load()
+            group_end = (slot.subview(5)).load()
+            group_idx = (slot.subview(7)).load()
+            if elect_one:
+                nvvm.mbarrier_arrive(sched_empty_mbar_ptr.subview(sched_stage))
+            sched_stage += 1
+            if sched_stage == SCHED_STAGES:
+                sched_stage = cutlass.Int32(0)
+                sched_full_phase = sched_full_phase ^ 1
+
+            if is_valid != 0:
+                coord_m_tile = group_begin + tile_m * cgrp_tile_mnk[0] + m_rank * cta_tile_mnk[0]
+                coord_n_c = tile_n * cgrp_tile_mnk[1] + n_rank * pair_n_size
+                if cutlass.const_expr(epi_rows_per_mma_m == 64):
+                    coord_n_c = coord_n_c + (warp_idx // 2) * cols_per_acc_stage
+
+                acc_stage = tile_iter % acc_stages
+                if acc_stage == 0 and tile_iter != 0:
+                    acc_full_phase_bit = acc_full_phase_bit ^ 1
+
+                while not nvvm.mbarrier_try_wait_parity(
+                    acc_full_mbar_ptr.subview(acc_stage),
+                    acc_full_phase_bit,
+                    time_limit=10_000_000,
+                ):
+                    pass
+
+                if cutlass.const_expr(use_acc_overlap):
+                    acc_buf_parity = tile_iter % 2
+                    acc_base_col = base_col_id_root + acc_buf_parity * acc_stage_stride
+                else:
+                    acc_buf_parity = cutlass.Int32(0)
+                    acc_base_col = base_col_id_root + acc_stage * acc_region_cols
+                # The 2-CTA epilogue drains its own half of the instruction's M,
+                # epi_rows_per_mma_m rows at a time, so a CTA tile of num_mma_m blocks
+                # drains in num_mma_m passes over its own column region.
+                for mi in cutlass.range_constexpr(num_mma_m):
+                    coord_m = coord_m_tile + mi * epi_rows_per_mma_m
+                    mi_col_base = acc_base_col + mi * epi_cols_per_mma_m
+                    tmem_col_addr_gemms = [(row_id_with_warp_offset << 16) | (mi_col_base + g * acc_gemm_stride) for g in range(num_gemms)]
+
+                    if cutlass.const_expr(epi_rows_per_mma_m == 64):
+                        row = coord_m + (warp_idx % 2) * 32 + lane
+                        row_active = True
+                    else:
+                        row = coord_m + tidx
+                        row_active = True
+
+                    # @@INJECT_AUX_VIEWS@@
+
+                    for subtile_idx in cutlass.range(subtile_cnt, unroll_full=True):
+                        if cutlass.const_expr(use_acc_overlap):
+                            _sub = subtile_idx + (1 - acc_buf_parity) * (subtile_cnt - 1 - 2 * subtile_idx)
+                            subtile_col_offset = _sub * t2r_inst_repx
+                        else:
+                            subtile_col_offset = subtile_idx * t2r_inst_repx
+                        c_rmem_vecs = []
+                        for g in cutlass.range_constexpr(num_gemms):
+                            tmem = cutlass.inttoptr(
+                                tmem_col_addr_gemms[g] + subtile_col_offset,
+                                6,
+                                cutlass.Float32,
+                            )
+                            c_rmem_vecs.append(nvvm.tcgen05_ld(shape, tmem, num=t2r_inst_repx))
+                        c_rmem_vec = c_rmem_vecs[0]
+
+                        if use_acc_overlap and mi == num_mma_m - 1 and subtile_idx == acc_overlap_subtiles - 1:
+                            nvvm.tcgen05_wait(kind=nvvm.Tcgen05Wait.LOAD)
+                            nvvm.tcgen05_fence(nvvm.Tcgen05Fence.BEFORE_THREAD_SYNC)
+                            if elect_one:
+                                mbar_pair_ptr = nvvm.mapa(acc_empty_mbar_ptr.subview(acc_stage), pair_leader_rank)
+                                nvvm.mbarrier_arrive(mbar_pair_ptr, scope=nvvm.MemScope.CLUSTER, relaxed=True)
+
+                        col = coord_n_c + subtile_col_offset
+
+                        # @@STG_ONLY:BEGIN@@
+
+                        if row_active and row < group_end:
+                            for j in cutlass.range_constexpr(t2r_inst_repx // vsize):
+                                col_j = col + j * vsize
+                                if col_j + vsize <= N:
+                                    vec_f32 = c_rmem_vec[j * vsize : (j + 1) * vsize]
+
+                                    # @@INJECT_STG_VEC_BINDINGS@@
+
+                                    # @@INJECT_EPILOGUE@@
+
+                        # @@STG_ONLY:END@@
+
+                if cutlass.const_expr(not use_acc_overlap):
+                    nvvm.tcgen05_wait(kind=nvvm.Tcgen05Wait.LOAD)
+                    nvvm.tcgen05_fence(nvvm.Tcgen05Fence.BEFORE_THREAD_SYNC)
+                    if elect_one:
+                        mbar_pair_ptr = nvvm.mapa(acc_empty_mbar_ptr.subview(acc_stage), pair_leader_rank)
+                        nvvm.mbarrier_arrive(mbar_pair_ptr, scope=nvvm.MemScope.CLUSTER, relaxed=True)
+                tile_iter += 1
+
+        if cutlass.const_expr(use_acc_overlap):
+            nvvm.tcgen05_wait(kind=nvvm.Tcgen05Wait.LOAD)
+            nvvm.tcgen05_fence(nvvm.Tcgen05Fence.BEFORE_THREAD_SYNC)
+            if elect_one:
+                nvvm.mbarrier_arrive(tmem_dealloc_mbar_ptr)
+
+    if warp_idx == unused_warp_id:
+        nvvm.setmaxregister(prod_reg_count, nvvm.SetMaxRegisterAction.DECREASE)
+
+
+@cute.jit
+def _host(
+    problem_size: tuple,
+    first_token_offset: cute.Tensor,
+    a_tma_workspace: cute.Tensor,
+    # @@INJECT_HOST_AB_PARAMS@@
+    # @@INJECT_HOST_TAP_PARAMS@@
+    # @@INJECT_HOST_AUX_PARAMS@@
+    stream: _cuda.CUstream,
+) -> None:
+    # @@INJECT_HOST_AB_LISTS@@
+
+    m = problem_size[0]
+    n = problem_size[1]
+    k_sym = problem_size[2]
+    num_experts = problem_size[3]
+    num_groups = problem_size[4]
+    _stride_idx = 5
+    _a_stride_sets = []
+    for _ in cutlass.range_constexpr(num_a_operands):
+        _a_stride_sets.append(
+            (
+                problem_size[_stride_idx],
+                problem_size[_stride_idx + 1],
+                problem_size[_stride_idx + 2],
+            )
+        )
+        _stride_idx += 3
+    _b_stride_sets = []
+    for _ in cutlass.range_constexpr(num_b_operands):
+        _b_stride_sets.append(
+            (
+                problem_size[_stride_idx],
+                problem_size[_stride_idx + 1],
+                problem_size[_stride_idx + 2],
+            )
+        )
+        _stride_idx += 3
+
+    # @@INJECT_HOST_REDUCTION_STRIDES@@
+
+    tma_a_desc_list = []
+    for _a_idx, _a_op in enumerate(_a_operands):
+        a_stride_m, a_stride_k, a_stride_l = _a_stride_sets[_a_idx]
+        tma_a_desc_list.append(
+            _tma.create_tensor_map_tiled(
+                global_address=_a_op.iterator.toint(),
+                dtype=ab_tma_desc_dtype,
+                global_dims=[k_sym, m, 1],
+                global_strides=[
+                    a_stride_m * ab_dtype.width // 128,
+                    a_stride_l * ab_dtype.width // 128,
+                ],
+                box_dims=[cta_tile_mnk[2], cta_tile_mnk[0], 1],
+                swizzle=ab_tma_swizzle,
+                tma_format=ab_tma_format,
+            )
+        )
+    tma_b_desc_list = []
+    for _b_idx, _b_op in enumerate(_b_operands):
+        b_stride_n, b_stride_k, b_stride_l = _b_stride_sets[_b_idx]
+        if cutlass.const_expr(b_is_n_major):
+            tma_b_desc_list.append(
+                _tma.create_tensor_map_tiled(
+                    global_address=_b_op.iterator.toint(),
+                    dtype=ab_tma_desc_dtype,
+                    global_dims=[n, k_sym, num_experts],
+                    global_strides=[
+                        b_stride_k * ab_dtype.width // 128,
+                        b_stride_l * ab_dtype.width // 128,
+                    ],
+                    box_dims=[b_tma_group_elems, cta_tile_mnk[2], 1],
+                    swizzle=ab_tma_swizzle,
+                    tma_format=ab_tma_format,
+                )
+            )
+        else:
+            tma_b_desc_list.append(
+                _tma.create_tensor_map_tiled(
+                    global_address=_b_op.iterator.toint(),
+                    dtype=ab_tma_desc_dtype,
+                    global_dims=[k_sym, n, num_experts],
+                    global_strides=[
+                        b_stride_n * ab_dtype.width // 128,
+                        b_stride_l * ab_dtype.width // 128,
+                    ],
+                    box_dims=[cta_tile_mnk[2], cta_tile_mnk[1], 1],
+                    swizzle=ab_tma_swizzle,
+                    tma_format=ab_tma_format,
+                )
+            )
+    rest_k = ((k_sym // block_size) + 3) // 4
+    rest_m = (m + 127) // 128 + num_groups
+    rest_n = (n + 127) // 128
+    tma_sfa_desc_list = []
+    for _sfa_op in _sfa_operands:
+        sfa_fp16_tensor = cute.make_tensor(
+            cute.recast_ptr(_sfa_op.iterator, dtype=cutlass.Float16),
+            cute.make_layout(
+                (256, rest_k, rest_m, 1),
+                stride=(
+                    1,
+                    256,
+                    cute.assume(256 * rest_k, 8),
+                    cute.assume(256 * rest_k * rest_m, 8),
+                ),
+            ),
+        )
+        tma_sfa_desc_list.append(
+            _tma.create_tensor_map_tiled_from_view(
+                sfa_fp16_tensor,
+                dtype=cutlass.Uint16,
+                box_dims=(256, sf_tma_box_k, sfa_tma_box_mn, 1),
+                stride_order=(0, 1, 2, 3),
+                swizzle=_tma.TensorMapSwizzle.none,
+            )
+        )
+    tma_sfb_desc_list = []
+    for _sfb_op in _sfb_operands:
+        sfb_fp16_tensor = cute.make_tensor(
+            cute.recast_ptr(_sfb_op.iterator, dtype=cutlass.Float16),
+            cute.make_layout(
+                (256, rest_k, rest_n, num_experts),
+                stride=(
+                    1,
+                    256,
+                    cute.assume(256 * rest_k, 8),
+                    cute.assume(256 * rest_k * rest_n, 8),
+                ),
+            ),
+        )
+        tma_sfb_desc_list.append(
+            _tma.create_tensor_map_tiled_from_view(
+                sfb_fp16_tensor,
+                dtype=cutlass.Uint16,
+                box_dims=(256, sf_tma_box_k, sfb_tma_box_mn, 1),
+                stride_order=(0, 1, 2, 3),
+                swizzle=_tma.TensorMapSwizzle.none,
+            )
+        )
+
+    cluster_m = cluster_shape_mnk[0]
+    cluster_n = cluster_shape_mnk[1]
+    grid_shape = (grid_num_clusters * cluster_m, cluster_n, 1)
+    _kernel(
+        problem_size[0],
+        problem_size[1],
+        problem_size[2],
+        cutlass.Int32(num_experts),
+        cutlass.Int32(num_groups),
+        first_token_offset,
+        a_tma_workspace,
+        # @@INJECT_HOST_KERNEL_DESC_PASS@@
+        # @@INJECT_MOE_HOST_MA_PASS@@
+        # @@INJECT_HOST_TAP_PASS@@
+        # @@INJECT_HOST_REDUCTION_STRIDE_PASS@@
+        # @@INJECT_HOST_AUX_PASS@@
+    ).launch(
+        grid=grid_shape,
+        block=(threads_per_cta, 1, 1),
+        cluster=cluster_shape_mnk,
+        use_pdl=USE_PDL,
+        stream=stream,
+    )
+
+
+@lru_cache(maxsize=None)
+def compile() -> Callable:
+    out_vec_elems = vec_bytes_epi // (cd_dtype.width // 8)
+    ab_stride_elems = 16 // (ab_dtype.width // 8)
+    sym_m = cute.sym_int64()
+    sym_n = cute.sym_int64(divisibility=out_vec_elems)
+    # K tails are supported: the K loop is ceil_div and the TMA descriptor's global K
+    # extent makes a partial box HW zero-filled. The only real K rule is the 16-byte
+    # TMA contiguous-extent one, already gated by _tma_alignment_reject.
+    sym_k = cute.sym_int64()
+    # Packed K extent: same reasoning as sym_k -- no CTA-tile multiple is required.
+    sym_kp = cute.sym_int64()
+    sym_e = cute.sym_int64()
+    sym_g = cute.sym_int64()
+
+    def _make_fake_a():
+        return make_fake_compact_tensor(
+            a_fake_dtype,
+            (sym_m, sym_kp, 1),
+            stride_order=(1, 0, 2),
+            assumed_align=16,
+        )
+
+    def _make_fake_b():
+        return make_fake_compact_tensor(
+            b_fake_dtype,
+            (sym_n, sym_kp, sym_e),
+            stride_order=(0, 1, 2) if b_is_n_major else (1, 0, 2),
+            assumed_align=16,
+        )
+
+    # SF reaches the kernel as a base pointer only; the host rebuilds the
+    # F8_128x4 view from problem_size. Modes 0/1 and all strides carry no
+    # contract; mode 2 keeps its literal plane count (1 for sfa, sym_e for sfb).
+    def _make_fake_sfa():
+        return cute.runtime.make_fake_tensor(
+            sf_cutlass_dtype,
+            (cute.sym_int64(), cute.sym_int64(), 1),
+            stride=(cute.sym_int64(), cute.sym_int64(), cute.sym_int64()),
+            assumed_align=16,
+        )
+
+    def _make_fake_sfb():
+        return cute.runtime.make_fake_tensor(
+            sf_cutlass_dtype,
+            (cute.sym_int64(), cute.sym_int64(), sym_e),
+            stride=(cute.sym_int64(), cute.sym_int64(), cute.sym_int64()),
+            assumed_align=16,
+        )
+
+    fake_first_token_offset = make_fake_compact_tensor(
+        offset_cutlass_dtype,
+        (sym_g,),
+        stride_order=(0,),
+        assumed_align=offset_cutlass_dtype.width // 8,
+    )
+    cluster_m = cluster_shape_mnk[0]
+    cluster_n = cluster_shape_mnk[1]
+    grid_ctas = grid_num_clusters * cluster_m * cluster_n
+    fake_a_tma_workspace = make_fake_compact_tensor(
+        cutlass.Int64,
+        (grid_ctas * num_a_operands * 16,),
+        stride_order=(0,),
+        assumed_align=128,
+    )
+
+    def _sym_operand_strides(is_mn_major: bool) -> tuple:
+        # Operand is permuted to (M|N, K, L): the unit stride is mode 0 when MN-major, mode 1 when K-major, and never reaches TMA.
+        unit = 0 if is_mn_major else 1
+        return tuple(cute.sym_int64() if i == unit else cute.sym_int64(divisibility=ab_stride_elems) for i in range(3))
+
+    sym_a_strides = []
+    for _ in range(num_a_operands):
+        sym_a_strides.extend(_sym_operand_strides(a_is_m_major))
+    sym_b_strides = []
+    for _ in range(num_b_operands):
+        sym_b_strides.extend(_sym_operand_strides(b_is_n_major))
+
+    # @@INJECT_COMPILE_REDUCTION_STRIDE_DECLS@@
+
+    # @@INJECT_COMPILE_AB_FAKES@@
+
+    # @@INJECT_COMPILE_TAP_FAKES@@
+
+    problem_size = (
+        sym_m,
+        sym_n,
+        sym_k,
+        sym_e,
+        sym_g,
+        *sym_a_strides,
+        *sym_b_strides,
+        # @@INJECT_COMPILE_REDUCTION_STRIDE_SYMBOLS@@
+    )
+
+    # @@INJECT_COMPILE_AUX_FAKES@@
+
+    _fake_stream = make_fake_stream(use_tvm_ffi_env_stream=False)
+    return cute.compile(
+        _host,
+        problem_size,
+        fake_first_token_offset,
+        fake_a_tma_workspace,
+        # @@INJECT_COMPILE_AB_PASS@@
+        # @@INJECT_COMPILE_TAP_PASS@@
+        # @@INJECT_COMPILE_AUX_PASS@@
+        stream=_fake_stream,
+        options=frost_compile_options,
+    )
diff --git a/python/cudnn/gemm/frost/tile_config.py b/python/cudnn/gemm/frost/tile_config.py
index 263d8e174..4901b0055 100644
--- a/python/cudnn/gemm/frost/tile_config.py
+++ b/python/cudnn/gemm/frost/tile_config.py
@@ -22,14 +22,16 @@
 
 @functools.lru_cache(maxsize=None)
 def _sm_smem_budget_bytes_of(device: int) -> int:
-    from cudnn.frost.device import device_name, is_available, shared_memory_per_block_optin
+    """Largest per-CTA SMEM the device gives a CTA — the oversized carveout where the
+    part has one, else the opt-in limit."""
+    from cudnn.frost.device import device_name, is_available, oversized_shared_memory_per_block, shared_memory_per_block_optin
 
     if not is_available():
         raise RuntimeError("cannot size the SMEM pipeline: no CUDA device is visible to query MaxSharedMemoryPerBlockOptin")
     optin = shared_memory_per_block_optin(device)
     if not optin:
         raise RuntimeError(f"the driver did not report MaxSharedMemoryPerBlockOptin for device {device_name(device)!r}; cannot size the SMEM pipeline")
-    return int(optin)
+    return max(int(optin), oversized_shared_memory_per_block(device))
 
 
 def _sm_smem_budget_bytes(device=None) -> int:
@@ -40,13 +42,11 @@ def _sm_smem_budget_bytes(device=None) -> int:
 
 # Per-CTA SMEM held back off the top when sizing the ab/acc pipeline, keyed by
 # the kernel template's pipeline: the CLC ring, smem barriers, the TMEM base-address
-_SMEM_FIXED_RESERVE_BY_PIPELINE = {"sm100": 2048, "sm103": 2048}
+_SMEM_FIXED_RESERVE_BY_PIPELINE = {"sm100": 2048, "sm103": 2048, "sm107": 2048}
 
 
 def _sm_smem_ab_budget_bytes(pipeline: str, device=None) -> int:
-    if pipeline not in _SMEM_FIXED_RESERVE_BY_PIPELINE:
-        raise NotImplementedError(f"SMEM fixed reserve not known for pipeline {pipeline!r}")
-    return _sm_smem_budget_bytes(device) - _SMEM_FIXED_RESERVE_BY_PIPELINE[pipeline]
+    return _sm_smem_budget_bytes(device) - _pipeline_fact(_SMEM_FIXED_RESERVE_BY_PIPELINE, pipeline, "SMEM fixed reserve")
 
 
 _L2_RETENTION_DIVISOR = 3
@@ -73,12 +73,22 @@ def l2_swizzle_budget_bytes(device=None) -> int:
 
 _CTA_TILE_M_MAX = 128
 _CTA_TILE_N_MAX = 256
-_CTA_TILE_K_BYTES_MAX = 128  # SWIZZLE_128B: SMEM row width = 128 bytes
 _MAX_CLUSTER_SIZE = _FROST_MAX_CLUSTER_SIZE
-_CTA_TILE_K_BYTES_MAX_BY_PIPELINE = {"sm100": 128, "sm103": 384}
+_CTA_TILE_K_BYTES_MAX_BY_PIPELINE = {"sm100": 128, "sm103": 384, "sm107": 128}
 
 # MMA-inst K in bytes
 _MMA_INST_K_BYTES = 32
+_MMA_INST_K_BYTES_BY_PIPELINE = {"sm100": 32, "sm103": 48, "sm107": 64}
+
+
+def _pipeline_fact(table: dict, pipeline: str, what: str):
+    """A per-pipeline hardware fact, by EXPLICIT membership. Every table keyed by
+    pipeline goes through here so a family added without an entry raises instead
+    of silently inheriting another family's value."""
+    if pipeline not in table:
+        raise NotImplementedError(f"{what} not known for pipeline {pipeline!r}; known: {sorted(table)}")
+    return table[pipeline]
+
 
 # MMA instructions the CTA tile spans along M. N is deliberately NOT an axis:
 # it measured within noise of a single instruction (0.96x vs 0.97x cuBLAS at
@@ -87,7 +97,7 @@ def l2_swizzle_budget_bytes(device=None) -> int:
 # words are indexed per 128-column block).
 _NUM_MMA_MAX = 2
 
-_AB_STAGES_CAP = 8  # cap even if SMEM permits more
+_AB_STAGES_CAP = 16  # cap even if SMEM permits more
 
 
 def smem_max_ab_stages(
@@ -154,19 +164,21 @@ def __post_init__(self) -> None:
         m, n, kb = self.cta_tile_m, self.cta_tile_n, self.cta_tile_k_bytes
         cm, cn = self.cgrp_size_m, self.cgrp_size_n
 
-        kb_max = _CTA_TILE_K_BYTES_MAX_BY_PIPELINE.get(self.pipeline, _CTA_TILE_K_BYTES_MAX)
-        if kb <= 0 or kb > kb_max or kb % _MMA_INST_K_BYTES != 0:
+        kb_max = _pipeline_fact(_CTA_TILE_K_BYTES_MAX_BY_PIPELINE, self.pipeline, "max cta_tile_k_bytes")
+        if kb <= 0 or kb > kb_max:
+            raise NotImplementedError(f"TileConfig {self.name!r}: cta_tile_k_bytes={kb} — must be " f"positive, ≤ {kb_max} for pipeline {self.pipeline}")
+        # sm103's K-tile is not free geometry either (K-tile = lcm(128, 48)).
+        if self.pipeline == "sm103" and kb != 384:
+            raise NotImplementedError(f"TileConfig {self.name!r}: sm103 fixes cta_tile_k_bytes=384 " f"(K-tile = lcm(128, 48)); got {kb}")
+        # A pipeline whose MMA instruction fixes its K width owns that axis — it
+        # is not free geometry (sm103 K=48B UTCOMMA, sm107 K=64B), and the K-tile
+        # walks that instruction, so it is a multiple of the SAME width.
+        mkb_want = _pipeline_fact(_MMA_INST_K_BYTES_BY_PIPELINE, self.pipeline, "MMA-inst K width")
+        if self.mma_inst_k_bytes != mkb_want:
+            raise NotImplementedError(f"TileConfig {self.name!r}: {self.pipeline} fixes " f"mma_inst_k_bytes={mkb_want}; got {self.mma_inst_k_bytes}")
+        if kb % mkb_want != 0:
             raise NotImplementedError(
-                f"TileConfig {self.name!r}: cta_tile_k_bytes={kb} — must be "
-                f"a positive multiple of {_MMA_INST_K_BYTES}, ≤ {kb_max} for "
-                f"pipeline {self.pipeline}"
-            )
-        # sm103 K axes are NOT free geometry
-        if self.pipeline == "sm103" and (kb != 384 or self.mma_inst_k_bytes != 48):
-            raise NotImplementedError(
-                f"TileConfig {self.name!r}: sm103 fixes cta_tile_k_bytes=384 and "
-                f"mma_inst_k_bytes=48 (K=48B UTCOMMA); got kb={kb}, "
-                f"mma_inst_k_bytes={self.mma_inst_k_bytes}"
+                f"TileConfig {self.name!r}: cta_tile_k_bytes={kb} — must be " f"a multiple of {self.pipeline}'s mma_inst_k_bytes={mkb_want}"
             )
 
         # CGRP size sanity. (cta_group-specific constraints — e.g. cgrp_size_m
@@ -348,9 +360,17 @@ class ConfigSm103(TileConfig):
     axes match :class:`ConfigSm100`. Type marker; callers pass ``pipeline="sm103"``."""
 
 
+class ConfigSm107(TileConfig):
+    """sm107 geometry — identical to :class:`ConfigSm100` except
+    ``mma_inst_k_bytes`` is 64 (the SM 10.7 block-scale MMA reads a 64-byte K
+    per instruction, twice sm100's 32). Type marker; callers pass
+    ``pipeline="sm107"``."""
+
+
 _CONFIG_CLASS_BY_PIPELINE: dict[str, type[TileConfig]] = {
     "sm100": ConfigSm100,
     "sm103": ConfigSm103,
+    "sm107": ConfigSm107,
 }
 
 
@@ -424,6 +444,22 @@ def _geom_sm103(cta_m: int, cta_n: int, cgrp_m: int, cgrp_n: int) -> ConfigSm103
     )
 
 
+def _geom_sm107(cta_m: int, cta_n: int, cgrp_m: int, cgrp_n: int) -> ConfigSm107:
+    """Build one sm107 config (the 64-byte MMA-inst K is the family's, not ours)."""
+    return ConfigSm107(
+        cta_tile_m=cta_m,
+        cta_tile_n=cta_n,
+        cta_tile_k_bytes=128,
+        cgrp_size_m=cgrp_m,
+        cgrp_size_n=cgrp_n,
+        epi_tile_mn=(cta_m, 32),
+        threads_per_cta=256,
+        pipeline="sm107",
+        acc_stages=2,
+        mma_inst_k_bytes=64,
+    )
+
+
 def _build_catalog() -> tuple[TileConfig, ...]:
     cfgs: list[TileConfig] = []
     for cta_m in (128, 64):
@@ -437,6 +473,12 @@ def _build_catalog() -> tuple[TileConfig, ...]:
     for cta_n in (256, 128):
         for cgrp_m, cgrp_n in _CLUSTERS:
             cfgs.append(_geom_sm103(128, cta_n, cgrp_m, cgrp_n))
+    # sm107 block-scale geometries: the sm100 axes narrowed to what the F8_128x4
+    # SF swizzle admits (M/N multiples of 128, K-tile 128 B) — the rest of the
+    # sm100 enumeration would only be rejected by validate_block_scale_config.
+    for cta_n in (256, 128):
+        for cgrp_m, cgrp_n in _CLUSTERS:
+            cfgs.append(_geom_sm107(128, cta_n, cgrp_m, cgrp_n))
     return tuple(cfgs)
 
 
@@ -591,7 +633,9 @@ def select_config(
     applied to the scored result, so they hold exactly as before.
 
     ``K`` is optional only for callers that do not have it to hand; without it the
-    small-K bias is neutral and everything else is unchanged.
+    small-K bias is neutral and everything else is unchanged. The pick is an sm100
+    geometry; a caller building for another template family passes it through
+    :func:`as_pipeline`.
     """
     sm = sm_count if sm_count is not None else _sm_count()
     x = max(1, num_gemms)
@@ -657,10 +701,35 @@ def select_config(
     # fusion, MoE and multi-GEMM have none and would fail template lookup.
     scheduler = "static" if (supports_static and x == 1 and M > 128 and not block_scale) else "clc"
 
-    name = f"CONFIG_sm100_{cta_m}x{cta_n}x128_{cta_m}x{cta_n}x32_cluster{cgrp_m}x{cgrp_n}"
+    name = f"CONFIG_sm100_{cta_m}x{cta_n}x128_{cta_m}x{cta_n}x{_MMA_INST_K_BYTES}_cluster{cgrp_m}x{cgrp_n}"
     return by_name(name), cta_group, scheduler
 
 
+def as_pipeline(cfg: TileConfig, pipeline: str) -> TileConfig:
+    """The same geometry as a ``pipeline``-family config — only the family-fixed
+    MMA-inst K width moves. A family whose K axes this geometry cannot satisfy
+    (sm103 fixes a 384-byte K-tile) raises from the config's ``__post_init__``,
+    so the invariant stays in one place."""
+    if cfg.pipeline == pipeline:
+        return cfg
+    cls = config_class_for_pipeline(pipeline)
+    return cls(
+        cta_tile_m=cfg.cta_tile_m,
+        cta_tile_n=cfg.cta_tile_n,
+        cta_tile_k_bytes=cfg.cta_tile_k_bytes,
+        cgrp_size_m=cfg.cgrp_size_m,
+        cgrp_size_n=cfg.cgrp_size_n,
+        epi_tile_mn=cfg.epi_tile_mn,
+        threads_per_cta=cfg.threads_per_cta,
+        pipeline=pipeline,
+        acc_stages=cfg.acc_stages,
+        tile_swizzle_n=cfg.tile_swizzle_n,
+        mma_inst_m=cfg.mma_inst_m,
+        mma_inst_n=cfg.mma_inst_n,
+        mma_inst_k_bytes=_pipeline_fact(_MMA_INST_K_BYTES_BY_PIPELINE, pipeline, "MMA-inst K width"),
+    )
+
+
 # ---------------------------------------------------------------------------
 # Block-scaled matmul config validation (geometry-only; cta_group lives on the
 # template). The F8_128x4 SF swizzle + 32x128b.warpx4 utccp atom impose:
diff --git a/python/properties.cpp b/python/properties.cpp
index f1be8b8db..552731ab1 100644
--- a/python/properties.cpp
+++ b/python/properties.cpp
@@ -143,6 +143,7 @@ init_properties(py::module_& m) {
         .value("FP8_E5M2", cudnn_frontend::DataType_t::FP8_E5M2)
         .value("FAST_FLOAT_FOR_FP8", cudnn_frontend::DataType_t::FAST_FLOAT_FOR_FP8)
         .value("FP8_E8M0", cudnn_frontend::DataType_t::FP8_E8M0)
+        .value("FP8_E5M3", cudnn_frontend::DataType_t::FP8_E5M3)
         .value("FP4_E2M1", cudnn_frontend::DataType_t::FP4_E2M1)
         .value("INT4", cudnn_frontend::DataType_t::INT4)
         .value("NOT_SET", cudnn_frontend::DataType_t::NOT_SET);
diff --git a/test/python/gemm/frost/gemm_test_utils.py b/test/python/gemm/frost/gemm_test_utils.py
index 16c285476..ebe53420d 100644
--- a/test/python/gemm/frost/gemm_test_utils.py
+++ b/test/python/gemm/frost/gemm_test_utils.py
@@ -26,6 +26,13 @@ def _active_sm() -> int | None:
 
 _SM = _active_sm()
 
+
+def _int8_mma_arch_ranges() -> tuple[tuple[int, int], ...]:
+    from cudnn.gemm.frost.kernel_registry import MMA_GPU_ARCH_SPECIAL_CASES
+
+    return MMA_GPU_ARCH_SPECIAL_CASES[("sm100", ("int8", "int8", "int32"))]
+
+
 # Every e2e test in this suite JITs sm100-family templates, valid only on
 # 100 <= SM < 120 (see kernel_registry.PIPELINE_ARCH_RANGES) — gate on arch, not just
 # GPU presence, so wrong-arch machines skip instead of failing in the JIT.
@@ -34,6 +41,22 @@ def _active_sm() -> int | None:
     reason="needs a Blackwell-family GPU (100 <= SM < 120), have " + ("none" if _SM is None else f"sm_{_SM}"),
 )
 
+# The int8 tcgen05 MMA is narrower than its family — SM 10.7 has no such
+# instruction, and NVVM fails to lower it rather than the JIT rejecting it.
+# Read the ranges off the registry so the suite never holds a second copy.
+INT8_SM_RANGES = _int8_mma_arch_ranges()
+requires_int8_mma = pytest.mark.skipif(
+    _SM is None or not any(lo <= _SM < hi for lo, hi in INT8_SM_RANGES),
+    reason="int8 MMA exists only on " + " or ".join(f"{lo} <= SM < {hi}" for lo, hi in INT8_SM_RANGES) + ", have " + ("none" if _SM is None else f"sm_{_SM}"),
+)
+
+# The sm107 templates render anywhere (the 64-byte-K mode is an idesc field, and
+# the OMMA descriptor is a host-side bit-pack); they RUN only on 107 <= SM < 110.
+requires_sm107 = pytest.mark.skipif(
+    _SM is None or not (107 <= _SM < 110),
+    reason="sm107 kernels run only on 107 <= SM < 110, have " + ("none" if _SM is None else f"sm_{_SM}"),
+)
+
 
 # --- plan / config resolution ----------------------------------------------
 
@@ -158,12 +181,60 @@ def rand_e8m0(shape, dev):
     return torch.randint(125, 129, shape, dtype=torch.uint8, device=dev).view(torch.float8_e8m0fnu)
 
 
+def e5m3_to_float(b: torch.Tensor) -> torch.Tensor:
+    """Decode E5M3 bytes: unsigned, 5-bit exponent (bias 15), 3-bit mantissa.
+
+    Exact over the whole 8-bit domain, matching the epilogue's decode:
+    ``E >= 1`` is the normal ``2^(E-15) * (1 + M/8)``; ``E == 0`` is subnormal
+    ``2^-14 * M/8 == M * 2^-17`` (so byte 0 is 0.0). ``E == 31`` is NOT inf/NaN —
+    the format is canonical-NaN-only, so only byte 255 is NaN and 248..254 are
+    finite (up to 114688); byte 255's value here is unused, since the hardware's
+    satfinite cvt never emits it."""
+    E = (b >> 3).to(torch.float32)
+    M = (b & 7).to(torch.float32)
+    normal = torch.pow(2.0, E - 15.0) * (1.0 + M / 8.0)
+    return torch.where(E >= 1, normal, M * (2.0**-17))
+
+
+def e5m3_finite_values(dev="cpu") -> torch.Tensor:
+    """The 255 finite E5M3 values, indexed by byte. Monotonically increasing
+    (the format is unsigned), so byte 254 = 114688 is the max finite and 255 is
+    the canonical NaN the hardware's satfinite cvt never emits."""
+    b = torch.arange(255, dtype=torch.float32, device=dev)
+    return e5m3_to_float(b.to(torch.int32))
+
+
+def e5m3_quant_ref(x: torch.Tensor) -> torch.Tensor:
+    """fp32 -> E5M3 byte, rounding UP, saturating to max finite — the reference
+    for what the epilogue's `cvt.rp.satfinite.ue5m3x2.f32` emits. Negative
+    inputs cannot occur (a scale is |amax| / output_max)."""
+    vals = e5m3_finite_values(x.device)
+    return torch.searchsorted(vals, x.clamp(min=0.0).contiguous(), right=False).clamp(max=254).to(torch.uint8)
+
+
+def rand_e5m3(shape, dev):
+    """Random E5M3 scale factors as raw bytes (torch has no E5M3 dtype, and the
+    kernel takes the SF blob as a base pointer anyway). Exponent field 13..17
+    keeps a scale PAIR inside FP32 while every value stays exactly
+    representable, so the torch reference is exact."""
+    return torch.randint(13 * 8, 18 * 8, shape, dtype=torch.uint8, device=dev)
+
+
 def block_quant_ref(x, block_size, out_dtype, scale_dtype):
-    """Torch reference for the block-quant epilogue: per-block amax scale
-    (E8M0 scales round toward +inf) + quantized output."""
+    """Torch reference for the block-quant epilogue: per-block amax scale +
+    quantized output. ``scale_dtype`` is a torch dtype, or the string
+    ``"e5m3"`` — torch has no E5M3, so that scale comes back as raw BYTES
+    (which is also what the kernel writes and what the test compares).
+
+    E8M0 and E5M3 scales round toward +inf; E4M3 scales round to nearest."""
     blocks = x.view(1, x.shape[0], x.shape[1] // block_size, block_size)
     output_max = 448.0 if out_dtype is torch.float8_e4m3fn else 57344.0
     scale_f = blocks.abs().amax(dim=-1) / output_max
+    if scale_dtype == "e5m3":
+        scale = e5m3_quant_ref(scale_f)
+        inv = torch.where(e5m3_to_float(scale.to(torch.int32)) > 0, e5m3_to_float(scale.to(torch.int32)).reciprocal(), 0.0)
+        q = (blocks * inv.unsqueeze(-1)).clamp(-output_max, output_max)
+        return q.to(out_dtype).view(1, x.shape[0], x.shape[1]), scale
     if scale_dtype is torch.float8_e8m0fnu:
         safe = torch.where(scale_f > 0, scale_f, 1.0)
         scale_f = torch.where(scale_f > 0, torch.pow(2.0, torch.ceil(torch.log2(safe))), 0.0)
diff --git a/test/python/gemm/frost/test_block_scale_matmul.py b/test/python/gemm/frost/test_block_scale_matmul.py
index ad5734a12..f514e00f1 100644
--- a/test/python/gemm/frost/test_block_scale_matmul.py
+++ b/test/python/gemm/frost/test_block_scale_matmul.py
@@ -19,6 +19,7 @@
 from gemm_test_utils import (
     _SM,
     requires_sm100,
+    requires_sm107,
     Plan as _plan,
     vp_bs as _vp_bs,
     kw as _kw,
@@ -27,12 +28,15 @@
     to_blocked as _to_blocked,
     unpack_fp4 as _unpack_fp4,
     rand_e8m0 as _rand_e8m0,
+    rand_e5m3 as _rand_e5m3,
+    e5m3_to_float as _e5m3_to_float,
     block_quant_ref as _block_quant_ref,
     reduction_ref as _reduction_ref,
     assert_block_scale_reduction_close as _assert_block_scale_reduction_close,
 )
 
 from cudnn.gemm.frost import compiler as C
+from cudnn.gemm.frost.dtypes import DTYPE_FROM_CUDNN as _DTYPE_FROM_CUDNN
 from cudnn.gemm.frost.compiler import jit_from_cudnn_graph
 from cudnn.gemm.frost.graph_analyzer import analyze
 from cudnn.gemm.frost.kernel_registry import GraphType, TEMPLATES, select_template
@@ -40,6 +44,7 @@
     CATALOG,
     ConfigSm100,
     ConfigSm103,
+    ConfigSm107,
     TileConfig,
     by_name,
     validate_block_scale_config,
@@ -261,9 +266,9 @@ def test_block_scale_matmul_gate_rejects_mismatches():
     # Missing F8_128x4 SF reorder layout.
     with pytest.raises(NotImplementedError, match="does not support"):
         _check_block_scale_supported(analyze(_build_nvfp4_graph(256, 256, 512, block_size=16, sf_dt=_DT_E4M3, reorder=False)), "sm100")
-    # nvfp4 (fp4+e4m3) with block32 — no supported case.
+    # FP8 data at block 16 — the fp8 rows are block-32 only, on every pipeline.
     with pytest.raises(NotImplementedError, match="does not support"):
-        _check_block_scale_supported(analyze(_build_nvfp4_graph(256, 256, 512, block_size=32, sf_dt=_DT_E4M3)), "sm100")
+        _check_block_scale_supported(analyze(_build_nvfp4_graph(256, 256, 512, block_size=16, sf_dt=_DT_E8M0, a_dt=_DT_E4M3)), "sm100")
     # mixed FP4 A / FP8 B (cross-family) — unsupported.
     with pytest.raises(NotImplementedError, match="does not support"):
         _check_block_scale_supported(
@@ -302,7 +307,7 @@ def test_analyze_detects_nvfp4_block_scale_matmul():
     chain = analyze(_build_nvfp4_graph(128, 256, 256, block_size=16))
     assert chain.has_block_scale
     bs = chain.block_scale
-    assert bs.combo == "nvfp4"
+    assert bs.a_dtype == "fp4_e2m1"
     assert bs.block_size == 16
     assert bs.sf_dtype == "fp8_e4m3"
     assert bs.mma_block_scale_kind == "MXF4NVF4"
@@ -330,7 +335,7 @@ def test_analyze_detects_mxfp8_block_scale_matmul():
         )
     )
     bs = chain.block_scale
-    assert bs.combo == "mxfp8"
+    assert bs.a_dtype == "fp8_e4m3"
     assert bs.block_size == 32 and bs.sf_dtype == "fp8_e8m0"
     assert bs.mma_block_scale_kind == "MXF8F6F4"
     assert bs.scale_vec_size == "BLOCK32"
@@ -340,7 +345,7 @@ def test_analyze_detects_mxfp8_block_scale_matmul():
 def test_analyze_detects_mxfp4_block_scale_matmul():
     chain = analyze(_build_nvfp4_graph(128, 256, 256, block_size=32, sf_dt=cudnn.data_type.FP8_E8M0))
     bs = chain.block_scale
-    assert bs.combo == "mxfp4"
+    assert bs.a_dtype == "fp4_e2m1"
     assert bs.block_size == 32 and bs.sf_dtype == "fp8_e8m0"
     assert bs.mma_block_scale_kind == "MXF4NVF4"
 
@@ -430,7 +435,8 @@ def _run_bs_numeric(combo, config_name, M, N, K, out_major="n"):
 
     g = _build_nvfp4_graph(M, N, K, block_size=bs, sf_dt=sf_dt, a_dt=a_dt, out_major=out_major)
     compiled = _plan(g, **_kw(config_name))
-    assert compiled.block_scale and compiled.chain.block_scale.combo == combo
+    assert compiled.block_scale
+    assert (compiled.chain.block_scale.sf_dtype, compiled.chain.block_scale.block_size) == (_DTYPE_FROM_CUDNN[sf_dt], bs)
 
     if out_major == "m":
         c = torch.zeros(1, N, M, dtype=torch.float16, device=dev).transpose(1, 2)
@@ -767,10 +773,13 @@ def _run_bs_quant_numeric(
 
     q = torch.empty(1, M, N, dtype=out_torch_dt, device=dev)
     q_scale_shape = scale_dim if scale_dim is not None else (1, M, N // 32)
+    # torch has no E5M3 dtype; the kernel writes the scale through an int8 byte
+    # carrier (a raw_ptr store to a uint8 tensor is rejected by the DSL).
+    scale_buf_dt = torch.int8 if scale_torch_dt == "e5m3" else scale_torch_dt
     if scale_reorder:
-        q_scale = torch.zeros(*q_scale_shape, dtype=scale_torch_dt, device=dev)
+        q_scale = torch.zeros(*q_scale_shape, dtype=scale_buf_dt, device=dev)
     else:
-        q_scale = torch.empty(*q_scale_shape, dtype=scale_torch_dt, device=dev)
+        q_scale = torch.empty(*q_scale_shape, dtype=scale_buf_dt, device=dev)
     aux = () if global_scale_tensor is None else (global_scale_tensor,)
     sf_k_padded = _ceil_div(K // bs, 4) * 4
     sfa_rows_padded = _ceil_div(M, 128) * 128
@@ -790,8 +799,11 @@ def _run_bs_quant_numeric(
 
     q_ref, scale_ref = _block_quant_ref(ref, 32, out_torch_dt, scale_torch_dt)
     if scale_reorder:
-        scale_ref = _to_blocked(scale_ref[0]).view_as(q_scale)
-    torch.testing.assert_close(q_scale.float(), scale_ref.float(), atol=0, rtol=0)
+        scale_ref = _to_blocked(scale_ref[0]).view_as(q_scale.view(scale_ref.dtype))
+    # E5M3 scales are compared as raw BYTES — the strictest form, and the only
+    # one available since torch cannot interpret the format.
+    got_scale = q_scale.view(torch.uint8) if scale_torch_dt == "e5m3" else q_scale
+    torch.testing.assert_close(got_scale.float(), scale_ref.float(), atol=0, rtol=0)
     torch.testing.assert_close(q.float(), q_ref.float(), atol=0, rtol=0)
 
 
@@ -1337,6 +1349,11 @@ def test_config_families():
     # A raw-base construction can't bypass the family invariant either.
     with pytest.raises(NotImplementedError, match="fixes cta_tile_k_bytes=384"):
         TileConfig(cta_tile_k_bytes=128, pipeline="sm103", **kw)
+    # The K-tile walks the MMA instruction, so it is a multiple of the PIPELINE's
+    # K width: 96 B is three sm100 instructions but not a whole number of sm107's.
+    ConfigSm100(cta_tile_k_bytes=96, mma_inst_k_bytes=32, pipeline="sm100", **kw)
+    with pytest.raises(NotImplementedError, match="multiple of sm107's mma_inst_k_bytes=64"):
+        ConfigSm107(cta_tile_k_bytes=96, mma_inst_k_bytes=64, pipeline="sm107", **kw)
     # Catalog entries carry their family class (the template-pairing key).
     assert all(isinstance(c, ConfigSm103) for c in CATALOG if c.pipeline == "sm103")
     assert all(isinstance(c, ConfigSm100) for c in CATALOG if c.pipeline == "sm100")
@@ -1377,7 +1394,7 @@ def test_select_template_dispatches_on_config_arch():
     assert t103.file == "sm103_block_scale_matmul_1ctamma.py"
     from cudnn.gemm.frost.kernel_registry import PIPELINE_ARCH_RANGES
 
-    assert PIPELINE_ARCH_RANGES[t103.pipeline] == ((103, 104),)
+    assert PIPELINE_ARCH_RANGES[t103.pipeline] == ((103, 110),)
     t100 = select_template(chain, by_name("CONFIG_sm100_128x128x128_128x128x32_cluster1x1"), cta_group=1, scheduler="clc")
     assert t100.file == "sm100_block_scale_matmul_1ctamma.py"
     t103_2 = select_template(chain, by_name(_CFG_128), cta_group=2, scheduler="clc")
@@ -1441,10 +1458,35 @@ def test_mma_gpu_arch_special_cases(monkeypatch):
 
 def test_jit_rejects_wrong_active_arch(monkeypatch):
     monkeypatch.setattr(C, "_current_arch", lambda: 100)
-    with pytest.raises(NotImplementedError, match=r"103 <= SM < 104.*sm_100"):
+    with pytest.raises(NotImplementedError, match=r"103 <= SM < 110.*sm_100"):
         jit_from_cudnn_graph(_bs_chain(), **_sm103_kw(_CFG_128))
 
 
+def test_sm103_rejects_multi_mma_m():
+    """The sm103 chunk pipeline has NOT been adapted to a CTA tile spanning
+    several MMA instructions along M: it miscomputes (A reads unwritten SMEM in
+    K) and its ab_stages budget under-counts, so cta_tile_m=256 also overruns the
+    SMEM cap. Both are silent-wrong / launch-fail, so the template declines the
+    geometry outright. Drop `supports_multi_mma_m=False` when it is fixed."""
+    wide = by_name("CONFIG_sm103_256x128x384_128x128x48_cluster1x1")
+    assert wide.num_mma_m == 2
+    for t in (t for t in TEMPLATES if t.pipeline == "sm103"):
+        assert not t.supports_multi_mma_m
+        assert "num_mma_m=2" in t.multi_mma_m_reject(wide)
+        assert t.multi_mma_m_reject(by_name(_CFG_128)) is None
+    # The other pipelines DO implement it — the gate is sm103-specific.
+    for f in ("sm100_block_scale_matmul_1ctamma.py", "sm107_block_scale_matmul_1ctamma.py"):
+        t = next(t for t in TEMPLATES if t.file == f)
+        assert t.supports_multi_mma_m and t.multi_mma_m_reject(wide) is None
+
+
+@requires_sm103
+def test_sm103_multi_mma_m_is_declined_not_miscomputed():
+    """The gate reaches the JIT path, so the geometry raises instead of running."""
+    with pytest.raises(NotImplementedError, match="several MMA instructions along M"):
+        jit_from_cudnn_graph(_bs_chain(), **_sm103_kw("CONFIG_sm103_256x128x384_128x128x48_cluster1x1"))
+
+
 # Renderer (no GPU needed beyond graph build)
 
 
@@ -1545,16 +1587,16 @@ def test_render_rejects_mxfp8(_pretend_sm103):
         ("nvfp4", "CONFIG_sm103_128x128x384_128x128x48_cluster8x1", 1),
         ("nvfp4", "CONFIG_sm103_128x128x384_128x128x48_cluster4x2", 2),
         ("nvfp4", "CONFIG_sm103_128x128x384_128x128x48_cluster16x1", 2),
-        # CTA tile split across two MMA instructions along M. Not in the sm103
-        # catalog (cta_m=128 only) — reachable by `by_name` synthesis. RENDER
-        # ONLY: no SM 10.3 part is available to check numerics.
-        ("nvfp4", "CONFIG_sm103_256x128x384_128x128x48_cluster1x1", 1),
-        ("nvfp4", "CONFIG_sm103_256x128x384_128x128x48_cluster2x1", 2),
+        # num_mma_m > 1 used to be smoke-rendered here. It renders, but the
+        # numerics are wrong on silicon (measured on an SM 10.7 part, which the
+        # sm103 arch range covers), so the template now declines the geometry —
+        # see test_sm103_rejects_multi_mma_m.
     ],
 )
 def test_sm103_compile_smoke(_pretend_sm103, combo, config_name, cta_group):
     compiled = jit_from_cudnn_graph(_bs_chain(combo=combo), **_sm103_kw(config_name, cta_group))
-    assert compiled.chain.block_scale.combo == combo
+    _bs = compiled.chain.block_scale
+    assert (_bs.sf_dtype, _bs.block_size) == (("fp8_e4m3", 16) if combo == "nvfp4" else ("fp8_e8m0", 32))
 
 
 def test_tma_alignment_unified():
@@ -1620,7 +1662,8 @@ def _run_sm103_numeric(combo, config_name, M, N, K, cta_group=1):
 
     g = _build_nvfp4_graph(M, N, K, block_size=bs, sf_dt=sf_dt, a_dt=a_dt)
     compiled = _plan(g, **_sm103_kw(config_name, cta_group))
-    assert compiled.block_scale and compiled.chain.block_scale.combo == combo
+    assert compiled.block_scale
+    assert (compiled.chain.block_scale.sf_dtype, compiled.chain.block_scale.block_size) == (_DTYPE_FROM_CUDNN[sf_dt], bs)
 
     # The F8_128x4 reorder pads to 128-row × 4-SF blocks; view with the
     # padded dims (matters for M/N not multiples of 128).
@@ -1709,12 +1752,13 @@ def test_auto_config_is_accepted_by_the_registry(M, N):
     scale is the narrow case: the F8_128x4 SF swizzle needs 128-multiple tiles,
     so the plain 32/64 ladder is illegal. Pin the invariant the funnel owns:
     whatever the heuristic picks must be in ``candidates(chain)``."""
-    from cudnn.gemm.frost.kernel_registry import candidates
-    from cudnn.gemm.frost.tile_config import select_config
+    from cudnn.gemm.frost.kernel_registry import candidates, preferred_pipeline
+    from cudnn.gemm.frost.tile_config import as_pipeline, select_config
 
     chain = analyze(_build_nvfp4_graph(M, N, 512))
     assert chain.has_block_scale
     cfg, _cta_group, _sched = select_config(chain.matmul.M, chain.matmul.N, chain.num_gemms, block_scale=chain.has_block_scale)
+    cfg = as_pipeline(cfg, preferred_pipeline(chain))  # the config build_gemm_plan actually builds
     accepted = {c.name for _t, c in candidates(chain)}
     assert accepted, "the registry accepts no geometry at all for this chain"
     assert cfg.name in accepted, f"select_config picked {cfg.name!r}, which the registry rejects for this graph"
@@ -1783,3 +1827,573 @@ def test_sf_blob_must_be_packed(M, K, kind, rejected):
     else:
         compiled(vp)
         torch.cuda.synchronize()
+
+
+# ---------------------------------------------------------------------------
+# sm107 block-scale pipeline (the sm100 pipeline on the 64-byte-K MMA)
+# ---------------------------------------------------------------------------
+
+_SM107_128 = "CONFIG_sm107_128x128x128_128x128x64_cluster1x1"
+_SM107_256 = "CONFIG_sm107_128x256x128_128x256x64_cluster1x1"
+
+
+@pytest.fixture
+def _pretend_sm107(monkeypatch):
+    monkeypatch.setattr(C, "_current_arch", lambda: 107)
+
+
+def _sm107_kw(config_name, cta_group=1):
+    return dict(config=by_name(config_name), cta_group=cta_group, scheduler="clc")
+
+
+def test_catalog_has_sm107_geometries():
+    sm107 = [c for c in CATALOG if c.pipeline == "sm107"]
+    # 2 cta_n × the shared 15-cluster enumeration.
+    assert len(sm107) == 30
+    pat = re.compile(r"^CONFIG_sm107_128x(128|256)x128_128x(128|256)x64_cluster\d+x\d+$")
+    for c in sm107:
+        assert pat.match(c.name), c.name
+        assert isinstance(c, ConfigSm107)
+        assert c.cta_tile_m == 128 and c.cta_tile_k_bytes == 128
+        assert c.mma_inst_k_bytes == 64
+    assert by_name(_SM107_128).geometry_name == "128x128x128_128x128x64_cluster1x1"
+
+
+def test_sm107_family_fixes_the_mma_k_width():
+    """The 64-byte MMA K is the FAMILY's, not free geometry — while the rest of
+    the geometry axes stay exactly sm100's."""
+    kw = dict(
+        cta_tile_m=128,
+        cta_tile_n=128,
+        cta_tile_k_bytes=128,
+        cgrp_size_m=1,
+        cgrp_size_n=1,
+        epi_tile_mn=(128, 32),
+        threads_per_cta=256,
+        pipeline="sm107",
+    )
+    ConfigSm107(mma_inst_k_bytes=64, **kw)
+    with pytest.raises(NotImplementedError, match="sm107 fixes mma_inst_k_bytes=64"):
+        ConfigSm107(mma_inst_k_bytes=32, **kw)
+    # A raw-base construction can't bypass the family invariant either.
+    with pytest.raises(NotImplementedError, match="sm107 fixes mma_inst_k_bytes=64"):
+        TileConfig(**kw)
+
+
+def test_sm107_template_selection_and_arch_gate(_pretend_sm107):
+    chain = analyze(_bs_chain())
+    for cta_group, want in ((1, "sm107_block_scale_matmul_1ctamma.py"), (2, "sm107_block_scale_matmul_2ctamma.py")):
+        cfg = by_name(_SM107_128 if cta_group == 1 else "CONFIG_sm107_128x128x128_128x128x64_cluster2x1")
+        tmpl = select_template(chain, cfg, cta_group=cta_group, scheduler="clc")
+        assert tmpl.file == want
+        assert tmpl.accepts(chain, cfg) is None
+    # An sm100 config still pairs with the sm100 templates on the same GPU.
+    sm100_cfg = by_name("CONFIG_sm100_128x128x128_128x128x32_cluster1x1")
+    assert select_template(chain, sm100_cfg, cta_group=1, scheduler="clc").file == "sm100_block_scale_matmul_1ctamma.py"
+
+
+def test_sm107_templates_reject_older_blackwell(monkeypatch):
+    monkeypatch.setattr(C, "_current_arch", lambda: 100)
+    chain = analyze(_bs_chain())
+    tmpl = next(t for t in TEMPLATES if t.file == "sm107_block_scale_matmul_1ctamma.py")
+    assert "107 <= SM < 110" in tmpl.accepts(chain, by_name(_SM107_128))
+
+
+@pytest.mark.parametrize(
+    "combo,cta_n,omma,k_mode,scales_per_inst,word_atoms,idesc_dtype",
+    [
+        ("nvfp4", 128, True, 2, 8, 2, "cutlass.Float4E2M1FN"),
+        ("nvfp4", 256, True, 2, 8, 2, "cutlass.Float4E2M1FN"),
+        ("mxfp4", 128, True, 2, 4, 1, "cutlass.Float4E2M1FN"),
+        ("mxfp8", 128, False, 1, 2, 1, "cutlass.Float8E4M3FN"),
+    ],
+)
+def test_render_sm107_tile_constants(combo, cta_n, omma, k_mode, scales_per_inst, word_atoms, idesc_dtype):
+    """One MMA spans a 64-byte K, so it eats twice sm100's scales; when that
+    outgrows a 4-scale utccp atom the SF word spans word_atoms of them. fp4
+    rides the OMMA descriptor, mxfp8 the MX one; both take the real dtype."""
+    chain = analyze(_bs_chain(combo))
+    cfg = by_name(f"CONFIG_sm107_128x{cta_n}x128_128x{cta_n}x64_cluster1x1")
+    txt = C._render_block_scale_tile_constants(cfg, chain, 1)
+    got = dict(line.split(" = ", 1) for line in txt.splitlines() if " = " in line and not line.startswith("#"))
+    assert got["idesc_is_omma"] == str(omma)
+    assert got["mma_k_dim_mode"] == str(k_mode)
+    assert got["sf_scales_per_inst"] == str(scales_per_inst)
+    assert got["word_atoms"] == str(word_atoms)
+    assert got["idesc_a_dtype"] == idesc_dtype
+
+
+@pytest.mark.parametrize("pipeline,mma_k,runs_on_512", [("sm100", 32, True), ("sm107", 64, False)])
+def test_tmem_columns_follow_the_arch_not_the_pipeline(pipeline, mma_k, runs_on_512, monkeypatch):
+    """TMEM size is a property of the GPU, so an sm100-pipeline kernel gets SM
+    10.7's 576 columns just like an sm107 one — and past 512 the alloc has to
+    ask for the exclusive mode. The extra columns are what lets a 256-wide N
+    tile double-buffer its accumulator instead of overlapping the two.
+
+    Only the sm100 config is checked on a 512-column part: PIPELINE_ARCH_RANGES
+    confines sm107 kernels to 107..109, which always have 576, and this tile's
+    SFB span needs 520 of them."""
+    chain = analyze(_bs_chain())
+    cfg = by_name(f"CONFIG_{pipeline}_128x256x128_128x256x{mma_k}_cluster1x1")
+
+    def render():
+        txt = C._render_block_scale_tile_constants(cfg, chain, 1)
+        return dict(line.split(" = ", 1) for line in txt.splitlines() if " = " in line and not line.startswith("#"))
+
+    monkeypatch.setattr(C, "_current_arch", lambda: 107)
+    got = render()
+    assert got["num_tmem_alloc_cols"] == "576" and got["tmem_alloc_exclusive"] == "True"
+    assert got["acc_stages"] == "2" and got["use_acc_overlap"] == "False"
+
+    monkeypatch.setattr(C, "_current_arch", lambda: 100)
+    if not runs_on_512:
+        with pytest.raises(NotImplementedError, match="TMEM span reaches column 520"):
+            render()
+        return
+    got = render()
+    assert got["num_tmem_alloc_cols"] == "512" and got["tmem_alloc_exclusive"] == "False"
+    assert got["acc_stages"] == "1" and got["use_acc_overlap"] == "True"
+
+
+@requires_sm107
+@pytest.mark.parametrize("combo", ["nvfp4", "mxfp4", "mxfp8"])
+@pytest.mark.parametrize(
+    "config_name,cta_group",
+    [
+        (_SM107_128 + "_1ctamma", 1),
+        (_SM107_256 + "_1ctamma", 1),
+        ("CONFIG_sm107_128x128x128_128x128x64_cluster1x2_1ctamma", 1),
+        ("CONFIG_sm107_128x128x128_128x128x64_cluster2x1_2ctamma", 2),
+        ("CONFIG_sm107_128x256x128_128x256x64_cluster2x1_2ctamma", 2),
+        ("CONFIG_sm107_128x256x128_128x256x64_cluster2x2_2ctamma", 2),
+    ],
+    ids=lambda v: v if isinstance(v, str) else f"cta{v}",
+)
+def test_sm107_block_scale_matmul_numerics(combo, config_name, cta_group):
+    _run_bs_numeric(combo, config_name, 256, 256, 512)
+
+
+@requires_sm107
+@pytest.mark.parametrize("combo", ["nvfp4", "mxfp4", "mxfp8"])
+@pytest.mark.parametrize("cta_group", [1, 2])
+@pytest.mark.parametrize("cta_m,cta_n", [(128, 256), (256, 128), (256, 256)])
+def test_sm107_block_scale_matmul_multi_mma_m(combo, cta_group, cta_m, cta_n):
+    """The CTA tile spanning several MMA instructions along M, on the 64-byte-K
+    pipeline. This is where the two SF regions stop agreeing: at nvfp4 a scale
+    word spans word_atoms=2 atoms, and SFA is indexed per M block (one MMA
+    instruction covers one 128-row block, so its word must be contiguous) while
+    SFB is walked across all N blocks by one instruction. Both layouts collapse
+    to the same addresses at a single block, so only cta_m/cta_n = 256 tells
+    them apart -- 256x256 is the case where both regions split at once."""
+    cluster = "cluster1x1" if cta_group == 1 else "cluster2x1"
+    suffix = "1ctamma" if cta_group == 1 else "2ctamma"
+    geometry = f"CONFIG_sm107_{cta_m}x{cta_n}x128_128x{cta_n}x64_{cluster}"
+    assert by_name(geometry).num_mma_m == cta_m // 128
+    _run_bs_numeric(combo, f"{geometry}_{suffix}", 256, 256, 512)
+
+
+@requires_sm107
+@pytest.mark.parametrize(
+    "combo,config_name,M,N,K",
+    [
+        ("nvfp4", _SM107_128 + "_1ctamma", 256, 384, 768),  # multi-tile N
+        ("mxfp4", _SM107_256 + "_1ctamma", 256, 256, 4096),  # many K-tiles
+        ("mxfp8", "CONFIG_sm107_128x128x128_128x128x64_cluster2x1_2ctamma", 384, 512, 256),
+        ("nvfp4", "CONFIG_sm107_128x128x128_128x128x64_cluster4x1_2ctamma", 1024, 512, 512),
+        ("nvfp4", "CONFIG_sm107_128x128x128_128x128x64_cluster1x4_1ctamma", 512, 1024, 512),
+    ],
+    ids=lambda v: v if isinstance(v, str) else str(v),
+)
+def test_sm107_block_scale_matmul_shapes_and_clusters(combo, config_name, M, N, K):
+    _run_bs_numeric(combo, config_name, M, N, K)
+
+
+@requires_sm107
+@pytest.mark.parametrize(
+    "config_name",
+    [
+        _SM107_128 + "_1ctamma",
+        "CONFIG_sm107_128x256x128_128x256x64_cluster2x1_1ctamma",  # M-OOB with cgrp_m > 1
+        "CONFIG_sm107_128x128x128_128x128x64_cluster1x2_1ctamma",  # N tile/cluster > N
+        "CONFIG_sm107_128x256x128_128x256x64_cluster2x1_2ctamma",  # 2-CTA pair
+    ],
+)
+def test_sm107_nvfp4_oob_shape(config_name):
+    """M=23, N=56, K=736 — ceil-padded SF descriptors + M/N/K OOB, on the
+    64-byte-K MMA (the last K-tile is only 736 % 256 = 224 elements)."""
+    test_nvfp4_oob_shape(config_name)
+
+
+@requires_sm107
+@pytest.mark.parametrize(
+    "combo,config_name",
+    [
+        ("nvfp4", _SM107_128 + "_1ctamma"),
+        ("mxfp8", "CONFIG_sm107_128x256x128_128x256x64_cluster2x1_2ctamma"),
+    ],
+)
+def test_sm107_block_scale_matmul_m_major(combo, config_name):
+    _run_bs_numeric(combo, config_name, 256, 256, 512, out_major="m")
+
+
+@requires_sm107
+@pytest.mark.parametrize("scale_reorder", [False, True])
+@pytest.mark.parametrize("config_name", [_SM107_128 + "_1ctamma", "CONFIG_sm107_128x256x128_128x256x64_cluster2x1_2ctamma"])
+def test_e5m3_quant_epilogue(config_name, scale_reorder):
+    """The epilogue can PRODUCE E5M3 scales, bit-exact against the torch
+    reference. The `cvt ... ue5m3x2.f32` this needs exists only on sm_107 —
+    strictly narrower than CONSUMING E5M3 scales, where the format is a
+    descriptor field every pipeline emits."""
+    _run_bs_quant_numeric(
+        config_name,
+        256,
+        256,
+        512,
+        cudnn.data_type.FP8_E4M3,
+        torch.float8_e4m3fn,
+        cudnn.data_type.FP8_E5M3,
+        "e5m3",
+        scale_reorder=scale_reorder,
+    )
+
+
+def test_e5m3_quant_rejects_off_sm107(monkeypatch):
+    """...and off sm_107 it declines cleanly rather than emitting PTX ptxas
+    would reject."""
+    g = _build_block_scale_quant_graph(256, 256, 512, dequant_block_size=32, scale_dt=cudnn.data_type.FP8_E5M3)
+    for arch in (100, 103, 110, 120):
+        monkeypatch.setattr(C, "_current_arch", lambda a=arch: a)
+        with pytest.raises(NotImplementedError, match=f"sm_{arch}"):
+            jit_from_cudnn_graph(g, **_kw(_SM107_128 + "_1ctamma"))
+
+
+@requires_sm107
+@pytest.mark.parametrize("config_name", [_SM107_128 + "_1ctamma", "CONFIG_sm107_128x256x128_128x256x64_cluster2x1_2ctamma"])
+def test_sm107_block_scale_matmul_quant_epilogue(config_name):
+    _run_bs_quant_numeric(
+        config_name,
+        256,
+        256,
+        512,
+        cudnn.data_type.FP8_E4M3,
+        torch.float8_e4m3fn,
+        cudnn.data_type.FP8_E8M0,
+        torch.float8_e8m0fnu,
+    )
+
+
+@requires_sm107
+@pytest.mark.parametrize("mode", [cudnn.reduction_mode.ADD, cudnn.reduction_mode.AMAX], ids=("add", "amax"))
+def test_sm107_block_scale_matmul_reduction_scalar(mode):
+    _run_bs_reduction_numeric(
+        "nvfp4",
+        _SM107_128 + "_1ctamma",
+        128,
+        128,
+        256,
+        mode,
+        red_dims=[1, 1, 1],
+        red_stride=None,
+        ref_dims=(0, 1, 2),
+    )
+
+
+@requires_sm107
+@pytest.mark.parametrize("config_name", [_SM107_128 + "_1ctamma", "CONFIG_sm107_128x256x128_128x256x64_cluster2x1_2ctamma"])
+def test_sm107_mxfp8_m_major_a_n_major_b(config_name):
+    test_mxfp8_m_major_a_n_major_b(config_name, 256, 256, 512)
+
+
+def test_auto_path_prefers_sm107_where_it_has_a_template(monkeypatch):
+    """On SM 10.7 the auto path builds block-scale graphs with the sm107
+    pipeline; graph types sm107 has no template for (plain matmul, MoE) fall
+    back to sm100 on their own, and so does older Blackwell."""
+    from cudnn.gemm.frost.kernel_registry import preferred_pipeline
+    from cudnn.gemm.frost.tile_config import as_pipeline, select_config
+
+    pg = cudnn.pygraph(
+        io_data_type=cudnn.data_type.BFLOAT16,
+        intermediate_data_type=cudnn.data_type.FLOAT,
+        compute_data_type=cudnn.data_type.FLOAT,
+    )
+    pa = pg.tensor(name="A", dim=[1, 128, 128], stride=[128 * 128, 128, 1])
+    pb = pg.tensor(name="B", dim=[1, 128, 128], stride=[128 * 128, 1, 128])
+    pg.matmul(A=pa, B=pb, name="mm").set_output(True)
+
+    bs_chain = analyze(_bs_chain())
+    plain_chain = analyze(pg)
+
+    monkeypatch.setattr(C, "_current_arch", lambda: 107)
+    assert preferred_pipeline(bs_chain) == "sm107"
+    assert preferred_pipeline(plain_chain) == "sm100"
+
+    monkeypatch.setattr(C, "_current_arch", lambda: 100)
+    assert preferred_pipeline(bs_chain) == "sm100"
+
+    # select_config scores pure geometry (an sm100 config); moving it to another
+    # family touches only that family's fixed MMA-inst K.
+    geo = select_config(4096, 4096, 1, block_scale=True)[0]
+    assert as_pipeline(geo, "sm100") is geo
+    geo107 = as_pipeline(geo, "sm107")
+    assert geo107.pipeline == "sm107" and geo107.mma_inst_k_bytes == 64
+    assert (geo107.cta_tile_mn, geo107.cgrp_size_mn) == (geo.cta_tile_mn, geo.cgrp_size_mn)
+    assert geo107.cta_tile_k_bytes == geo.cta_tile_k_bytes
+    # sm103 fixes a 384-byte K-tile, so a scored geometry cannot become one — the
+    # family invariant lives on the config, not in a second whitelist.
+    with pytest.raises(NotImplementedError, match="cta_tile_k_bytes=384"):
+        as_pipeline(geo, "sm103")
+
+
+# ---------------------------------------------------------------------------
+# FP4 with E5M3 scale factors (SM 10.7+)
+#
+# UTCOMMA's instruction descriptor picks the SF format (0=E4M3, 1=E8M0,
+# 2=E5M3); only SM 10.7 decodes 2, and unlike nvfp4/mxfp4 -- which each fix one
+# K-block -- E5M3 is legal with both 16 and 32. Nothing else moves: the SF is
+# still a byte in the F8_128x4 layout, so the templates are untouched.
+# ---------------------------------------------------------------------------
+
+
+def _e5m3_graph(M, N, K, block_size, **kw):
+    return _build_nvfp4_graph(
+        M,
+        N,
+        K,
+        block_size=block_size,
+        sf_dt=cudnn.data_type.FP8_E5M3,
+        a_dt=cudnn.data_type.FP4_E2M1,
+        **kw,
+    )
+
+
+@pytest.mark.parametrize("block_size", [16, 32])
+def test_e5m3_analyzer_reads_the_scale_dtype(block_size):
+    bs = analyze(_e5m3_graph(256, 256, 512, block_size)).block_scale
+    assert bs.sf_dtype_a == "fp8_e5m3" and bs.sf_dtype_b == "fp8_e5m3"
+    assert bs.block_size == block_size
+    assert bs.a_dtype == "fp4_e2m1"
+    assert bs.sf_scale_format == 2
+
+
+@pytest.mark.parametrize(
+    "sf_dt,sf_name", [(cudnn.data_type.FP8_E4M3, "fp8_e4m3"), (cudnn.data_type.FP8_E8M0, "fp8_e8m0"), (cudnn.data_type.FP8_E5M3, "fp8_e5m3")]
+)
+@pytest.mark.parametrize("block_size", [16, 32])
+def test_fp4_scale_dtype_and_block_are_orthogonal(sf_dt, sf_name, block_size):
+    """All three FP4 scale-factor dtypes are legal at BOTH K-blocks — nvfp4
+    (e4m3/16) and mxfp4 (e8m0/32) are just the best-known corners, not a
+    coupling. The registry must carry the full 3x2 matrix on every pipeline that
+    has fp4 at all."""
+    from cudnn.gemm.frost.kernel_registry import GraphType as _GT, MMA_TYPE_SUPPORT, _bs_key
+
+    key = _bs_key("fp4_e2m1", sf_name, "fp4_e2m1", sf_name, block_size)
+    for pipeline in ("sm100", "sm103", "sm107"):
+        assert key in MMA_TYPE_SUPPORT[pipeline][_GT.BLOCK_SCALE_MATMUL], f"{pipeline} is missing fp4+{sf_name}/{block_size}"
+    # and the analyzer reads the pair back off a graph built with it
+    bs = analyze(_build_nvfp4_graph(256, 256, 512, block_size=block_size, sf_dt=sf_dt, a_dt=cudnn.data_type.FP4_E2M1)).block_scale
+    assert (bs.a_dtype, bs.sf_dtype, bs.block_size) == ("fp4_e2m1", sf_name, block_size)
+
+
+@_GPU
+@pytest.mark.parametrize("sf_dt,sf_name", [(cudnn.data_type.FP8_E4M3, "fp8_e4m3"), (cudnn.data_type.FP8_E8M0, "fp8_e8m0")])
+@pytest.mark.parametrize("block_size", [16, 32])
+@pytest.mark.parametrize("config_name", ["CONFIG_sm100_128x128x128_128x128x32_cluster1x1", _SM107_128])
+def test_fp4_all_scale_block_corners_numerics(config_name, sf_dt, sf_name, block_size):
+    """Numerics for the whole non-E5M3 fp4 matrix, including the two corners the
+    nvfp4 / mxfp4 pair leaves out: e4m3 at block 32 and e8m0 at block 16."""
+    dev = "cuda"
+    torch.manual_seed(0)
+    M, N, K = 256, 256, 512
+    sf_k = K // block_size
+    lut = torch.tensor(_E2M1, dtype=torch.float32, device=dev)
+    a_u8 = torch.randint(0, 256, (1, M, K // 2), dtype=torch.uint8, device=dev)
+    b_u8 = torch.randint(0, 256, (1, N, K // 2), dtype=torch.uint8, device=dev)
+    if sf_name == "fp8_e8m0":
+        sfa, sfb = _rand_e8m0((M, sf_k), dev), _rand_e8m0((N, sf_k), dev)
+    else:
+        sfa = torch.randint(1, 4, (M, sf_k), device=dev).to(torch.float8_e4m3fn)
+        sfb = torch.randint(1, 4, (N, sf_k), device=dev).to(torch.float8_e4m3fn)
+
+    g = _build_nvfp4_graph(M, N, K, block_size=block_size, sf_dt=sf_dt, a_dt=cudnn.data_type.FP4_E2M1)
+    compiled = _plan(g, config=by_name(config_name), cta_group=1, scheduler="clc")
+    assert (compiled.chain.block_scale.sf_dtype, compiled.chain.block_scale.block_size) == (sf_name, block_size)
+
+    c = torch.zeros(1, M, N, dtype=torch.float16, device=dev)
+    compiled(
+        _vp_bs(
+            compiled, a_u8.view(torch.float4_e2m1fn_x2), b_u8.view(torch.float4_e2m1fn_x2), c, _to_blocked(sfa).view(1, 1, -1), _to_blocked(sfb).view(1, 1, -1)
+        )
+    )
+    torch.cuda.synchronize()
+
+    a_s = _unpack_fp4(a_u8, lut).view(M, K) * sfa.float().repeat_interleave(block_size, 1)
+    b_s = _unpack_fp4(b_u8, lut).view(N, K) * sfb.float().repeat_interleave(block_size, 1)
+    torch.testing.assert_close(c[0], (a_s @ b_s.t()).to(torch.float16), atol=2e-1, rtol=2e-2)
+
+
+# Block-scale cases that only some GPUs decode: SM 10.7 added the E5M3 scale
+# format (either K-block) and E4M3 at block 32. Keyed by (SF dtype, K-block).
+_GPU_GATED_FP4_CASES = {("fp8_e5m3", 16), ("fp8_e5m3", 32), ("fp8_e4m3", 32)}
+_DTYPE_GATED_SF_DTYPES = {"fp8_e5m3"}
+_GPU_GATED_RANGES = ((107, 110),)
+
+
+def test_gpu_gated_cases_are_narrowed_everywhere():
+    """The load-bearing invariant behind putting the GPU-gated fp4 cases in the
+    ordinary case sets: EVERY one of them, on EVERY pipeline that carries it,
+    needs its own MMA_GPU_ARCH_SPECIAL_CASES entry. Miss one — a new K-block, a
+    new family inheriting _BLOCK_SCALE_CASES — and that combo is accepted on a
+    part whose descriptor cannot encode it: silently wrong scales, not a clean
+    rejection. The registry cannot derive this, so it is pinned here."""
+    from cudnn.gemm.frost.kernel_registry import GraphType as _GT, MMA_GPU_ARCH_SPECIAL_CASES, MMA_TYPE_SUPPORT, _bs_key
+
+    gated = [
+        (pipeline, _bs_key("fp4_e2m1", sf, "fp4_e2m1", sf, blk))
+        for pipeline, by_type in MMA_TYPE_SUPPORT.items()
+        for sf, blk in _GPU_GATED_FP4_CASES
+        if _bs_key("fp4_e2m1", sf, "fp4_e2m1", sf, blk) in by_type.get(_GT.BLOCK_SCALE_MATMUL, ())
+    ]
+    assert len(gated) == 3 * len(_GPU_GATED_FP4_CASES), f"expected every pipeline to carry every gated case, got {len(gated)}"
+    bad = [pk for pk in gated if MMA_GPU_ARCH_SPECIAL_CASES.get(pk) != _GPU_GATED_RANGES]
+    assert not bad, f"GPU-gated cases missing their {_GPU_GATED_RANGES} narrowing: {bad}"
+
+
+def test_dtype_and_mma_arch_gates_are_independent():
+    """A narrow DTYPE and a narrow MMA INSTRUCTION are separate facts that happen
+    to share a range today. Keep both: the dtype's range can widen on a later
+    part (E5M3 elsewhere than a block-scale MMA operand), while this fp4+E5M3
+    instruction's cannot. Collapsing either into the other would let one widen
+    the other silently."""
+    from cudnn.gemm.frost.dtypes import DTYPE_GPU_ARCH_RANGES
+    from cudnn.gemm.frost.kernel_registry import MMA_GPU_ARCH_SPECIAL_CASES, _bs_key
+
+    for sf in _DTYPE_GATED_SF_DTYPES:
+        assert DTYPE_GPU_ARCH_RANGES.get(sf) == _GPU_GATED_RANGES, f"{sf} is not narrowed by the dtype table"
+        for pipeline in ("sm100", "sm103", "sm107"):
+            for blk in (16, 32):
+                key = (pipeline, _bs_key("fp4_e2m1", sf, "fp4_e2m1", sf, blk))
+                assert MMA_GPU_ARCH_SPECIAL_CASES.get(key) == _GPU_GATED_RANGES, f"{key} lost its independent MMA-instruction narrowing"
+
+
+@pytest.mark.parametrize("sf_name,block_size", sorted(_GPU_GATED_FP4_CASES))
+@pytest.mark.parametrize("pipeline", ["sm100", "sm103", "sm107"])
+def test_gpu_gated_cases_reject_off_sm107(pipeline, sf_name, block_size, monkeypatch):
+    """...and the narrowing actually bites: accepted on 10.7/10.9, turned away
+    by the ARCH gate everywhere else, on every pipeline."""
+    from cudnn.gemm.frost.kernel_registry import GraphType as _GT, mma_arch_reject
+
+    sf_dt = {"fp8_e5m3": cudnn.data_type.FP8_E5M3, "fp8_e4m3": cudnn.data_type.FP8_E4M3}[sf_name]
+    chain = analyze(_build_nvfp4_graph(256, 256, 512, block_size=block_size, sf_dt=sf_dt, a_dt=cudnn.data_type.FP4_E2M1))
+    for arch in (107, 109):
+        monkeypatch.setattr(C, "_current_arch", lambda a=arch: a)
+        assert mma_arch_reject(chain, _GT.BLOCK_SCALE_MATMUL, pipeline) is None
+    for arch in (100, 103, 120):
+        monkeypatch.setattr(C, "_current_arch", lambda a=arch: a)
+        reason = mma_arch_reject(chain, _GT.BLOCK_SCALE_MATMUL, pipeline)
+        assert reason is not None, f"{pipeline} accepted fp4+{sf_name}/{block_size} on sm_{arch}"
+        assert f"sm_{arch}" in reason and "107 <= SM < 110" in reason, reason
+
+
+@pytest.mark.parametrize("block_size", [16, 32])
+def test_dtype_gated_scales_reject_off_sm107(block_size, monkeypatch):
+    """The dtype gate bites wherever the dtype is NAMED — it reads the chain, so
+    it does not care which pipeline, graph shape or code path would have run."""
+    from cudnn.gemm.frost.dtypes import dtype_arch_reject
+
+    chain = analyze(_build_nvfp4_graph(256, 256, 512, block_size=block_size, sf_dt=cudnn.data_type.FP8_E5M3, a_dt=cudnn.data_type.FP4_E2M1))
+    assert "fp8_e5m3" in chain.dtypes_used()
+    for arch in (107, 109):
+        assert dtype_arch_reject(chain, arch) is None
+    for arch in (100, 103, 110, 120):
+        reason = dtype_arch_reject(chain, arch)
+        assert reason is not None, f"accepted an E5M3 scale on sm_{arch}"
+        assert f"sm_{arch}" in reason and "107 <= SM < 110" in reason, reason
+    monkeypatch.setattr(C, "_current_arch", lambda: 100)
+    with pytest.raises(NotImplementedError, match="sm_100"):
+        jit_from_cudnn_graph(
+            _build_nvfp4_graph(256, 256, 512, block_size=block_size, sf_dt=cudnn.data_type.FP8_E5M3, a_dt=cudnn.data_type.FP4_E2M1),
+            **_kw(_SM107_128 + "_1ctamma"),
+        )
+
+
+@requires_sm107
+@pytest.mark.parametrize("block_size", [16, 32])
+def test_e5m3_runs_on_the_sm100_pipeline(block_size):
+    """The sm100 templates reach scale_format=2 through the MX descriptor rather
+    than the OMMA one, and SM 10.7 decodes it there too — so an sm100-pipeline
+    config is a legitimate way to run E5M3 on this part."""
+    _run_e5m3_numeric("CONFIG_sm100_128x128x128_128x128x32_cluster1x1", 1, block_size)
+
+
+@pytest.mark.parametrize("block_size,scales_per_inst,word_atoms", [(16, 8, 2), (32, 4, 1)])
+def test_render_e5m3_tile_constants(block_size, scales_per_inst, word_atoms):
+    """scale_format is the only constant E5M3 moves; the SF-word geometry still
+    follows block_size alone, exactly as for nvfp4 (16) and mxfp4 (32)."""
+    chain = analyze(_e5m3_graph(256, 256, 512, block_size))
+    txt = C._render_block_scale_tile_constants(by_name(_SM107_128), chain, 1)
+    got = dict(line.split(" = ", 1) for line in txt.splitlines() if " = " in line and not line.startswith("#"))
+    assert got["sf_scale_format"] == "2"
+    assert got["idesc_is_omma"] == "True"
+    assert got["sf_scales_per_inst"] == str(scales_per_inst)
+    assert got["word_atoms"] == str(word_atoms)
+
+
+def _run_e5m3_numeric(config_name, cta_group, block_size, M=256, N=256, K=512):
+    dev = "cuda"
+    torch.manual_seed(0)
+    sf_k = K // block_size
+    lut = torch.tensor(_E2M1, dtype=torch.float32, device=dev)
+    a_u8 = torch.randint(0, 256, (1, M, K // 2), dtype=torch.uint8, device=dev)
+    b_u8 = torch.randint(0, 256, (1, N, K // 2), dtype=torch.uint8, device=dev)
+    a_deq = _unpack_fp4(a_u8, lut).view(M, K)
+    b_deq = _unpack_fp4(b_u8, lut).view(N, K)
+    sfa, sfb = _rand_e5m3((M, sf_k), dev), _rand_e5m3((N, sf_k), dev)
+
+    g = _e5m3_graph(M, N, K, block_size)
+    compiled = _plan(g, config=by_name(config_name), cta_group=cta_group, scheduler="clc")
+    assert (compiled.chain.block_scale.sf_dtype, compiled.chain.block_scale.block_size) == ("fp8_e5m3", block_size)
+
+    # The SF blob is read by base pointer and to_blocked() ceil-pads to whole
+    # 128x4 atoms, so its element count is not M*sf_k for a ragged shape — pass
+    # the byte run itself rather than a logical view of it.
+    c = torch.zeros(1, M, N, dtype=torch.float16, device=dev)
+    compiled(
+        _vp_bs(
+            compiled,
+            a_u8.view(torch.float4_e2m1fn_x2),
+            b_u8.view(torch.float4_e2m1fn_x2),
+            c,
+            _to_blocked(sfa).view(1, 1, -1),
+            _to_blocked(sfb).view(1, 1, -1),
+        )
+    )
+    torch.cuda.synchronize()
+
+    a_s = a_deq * _e5m3_to_float(sfa).repeat_interleave(block_size, 1)
+    b_s = b_deq * _e5m3_to_float(sfb).repeat_interleave(block_size, 1)
+    torch.testing.assert_close(c[0], (a_s @ b_s.t()).to(torch.float16), atol=2e-1, rtol=2e-2)
+
+
+@requires_sm107
+@pytest.mark.parametrize("block_size", [16, 32])
+@pytest.mark.parametrize(
+    "config_name,cta_group",
+    [
+        (_SM107_128, 1),
+        (_SM107_256, 1),
+        ("CONFIG_sm107_128x128x128_128x128x64_cluster1x2", 1),
+        ("CONFIG_sm107_128x128x128_128x128x64_cluster2x1", 2),
+        ("CONFIG_sm107_128x256x128_128x256x64_cluster2x1", 2),
+    ],
+    ids=lambda v: v if isinstance(v, str) else f"cta{v}",
+)
+def test_e5m3_block_scale_matmul_numerics(config_name, cta_group, block_size):
+    _run_e5m3_numeric(config_name, cta_group, block_size)
+
+
+@requires_sm107
+@pytest.mark.parametrize("block_size", [16, 32])
+def test_e5m3_block_scale_matmul_oob_shape(block_size):
+    """M-OOB / K past a tile boundary is TMA zero-fill, same as every other
+    block-scale combo."""
+    _run_e5m3_numeric(_SM107_128, 1, block_size, M=255, N=256, K=512 + 4 * block_size)
diff --git a/test/python/gemm/frost/test_block_scale_matmul_swiglu.py b/test/python/gemm/frost/test_block_scale_matmul_swiglu.py
index c7c47fc11..27c2de1bb 100644
--- a/test/python/gemm/frost/test_block_scale_matmul_swiglu.py
+++ b/test/python/gemm/frost/test_block_scale_matmul_swiglu.py
@@ -10,6 +10,7 @@
 
 from gemm_test_utils import (
     requires_sm100,
+    requires_sm107,
     Plan as _plan,
     kw as _kw,
     E2M1 as _E2M1,
@@ -120,7 +121,7 @@ def test_shared_dequant_dedup():
     assert chain.num_b_operands == 2
     assert chain.gemm_operands == [(0, 0), (0, 1)]
     assert chain.has_block_scale
-    assert chain.block_scale.combo == "nvfp4"
+    assert (chain.block_scale.sf_dtype, chain.block_scale.block_size) == ("fp8_e4m3", 16)
 
 
 def test_shared_dequant_reduction_detected():
@@ -212,7 +213,8 @@ def _run(combo, config_name, M, N, K):
     g = _build_dual_bs_graph(M, N, K, combo=combo)
     compiled = _plan(g, **_kw(config_name))
     assert compiled.chain.is_multi_gemm and compiled.block_scale
-    assert compiled.chain.block_scale.combo == combo
+    _bs = compiled.chain.block_scale
+    assert (_bs.sf_dtype, _bs.block_size) == (("fp8_e4m3", 16) if combo == "nvfp4" else ("fp8_e8m0", 32))
 
     c = torch.zeros(1, M, N, dtype=torch.float32, device=dev)
     compiled(_vp_bs_mg(compiled, pairs, c))
@@ -500,3 +502,30 @@ def test_dual_block_scale_matmul_reduction_rejects_int32():
     )
     with pytest.raises(NotImplementedError, match="fp32 compute/output"):
         jit_from_cudnn_graph(g, **_kw("CONFIG_sm100_128x128x128_128x128x32_cluster1x1_1ctamma"))
+
+
+# ---------------------------------------------------------------------------
+# sm107 (64-byte-K MMA): the same dual block-scale SwiGLU chain
+# ---------------------------------------------------------------------------
+
+
+@requires_sm107
+@pytest.mark.parametrize("combo", ["nvfp4", "mxfp4", "mxfp8"])
+@pytest.mark.parametrize(
+    "config_name",
+    [
+        "CONFIG_sm107_128x128x128_128x128x64_cluster1x1_1ctamma",
+        "CONFIG_sm107_128x128x128_128x128x64_cluster2x1_2ctamma",
+    ],
+)
+def test_sm107_dual_block_scale_matmul_numerics(combo, config_name):
+    """Two parallel block-scale GEMMs sharing A + one epilogue. Both GEMMs land
+    in TMEM alongside the per-operand SF words — the tighter budget the 576
+    columns buy back."""
+    _run(combo, config_name, 256, 128, 512)
+
+
+@requires_sm107
+@pytest.mark.parametrize("combo", ["nvfp4", "mxfp8"])
+def test_sm107_dual_block_scale_matmul_swiglu_quant_epilogue(combo):
+    test_dual_block_scale_matmul_swiglu_quant_epilogue(combo, "CONFIG_sm107_128x128x128_128x128x64_cluster1x1_1ctamma")
diff --git a/test/python/gemm/frost/test_frontend_integration.py b/test/python/gemm/frost/test_frontend_integration.py
index a8a110655..d599cfaf4 100644
--- a/test/python/gemm/frost/test_frontend_integration.py
+++ b/test/python/gemm/frost/test_frontend_integration.py
@@ -4,8 +4,8 @@
 """Frontend integration: the GEMM engine ``frost_gemm`` joins the ONE ranked
 plan list the graph API builds at create_execution_plans() — discovered from
 ``cudnn/engines/manifest.py``, no registration call and no environment variable.
-Exercises its presence in the list, select_plan/deselect_engines, the build walk
-falling through a declining plan, ineligible graphs, and the wrapper.Graph path.
+Exercises its presence in the list, select_plan, ineligible graphs, and the
+wrapper.Graph path.
 """
 
 from __future__ import annotations
@@ -16,8 +16,7 @@
 from gemm_test_utils import requires_sm100
 
 import cudnn
-from cudnn.engines import MANIFEST, PlanConfig, heuristics, is_backend_engine, is_python_engine
-from test_dispatch import _FAKE, _offer
+from cudnn.engines import MANIFEST, is_backend_engine, is_python_engine
 
 pytestmark = pytest.mark.L0
 
@@ -230,11 +229,11 @@ def test_misaligned_buffer_rejected():
 
 
 @_GPU
-@pytest.mark.parametrize("route", ["default", "frost", "deselected"])
+@pytest.mark.parametrize("route", ["default", "frost"])
 def test_ranked_list_routes_and_all_routes_agree(route):
     """The same graph runs on whichever entry of the ranked list is selected:
-    the default, the FROST entry pinned by select_plan, or the backend after
-    deselect_engines bars FROST. All three produce the right numbers.
+    the default, or the FROST entry pinned by select_plan. Both produce the
+    right numbers.
 
     The default route asserts nothing about WHICH engine served it — that is
     the placeholder in engines/heuristics.py, and a cost model will change it."""
@@ -243,8 +242,6 @@ def test_ranked_list_routes_and_all_routes_agree(route):
     _plan(g)
     if route == "frost":
         _pin_frost(g)
-    elif route == "deselected":
-        g.deselect_engines([_FROST])
     g.check_support()
     g.build_plans()
     if route == "frost":
@@ -253,48 +250,6 @@ def test_ranked_list_routes_and_all_routes_agree(route):
         # scratch from the caller (its TMA-descriptor buffer is a one-time,
         # plan-owned allocation) — not a blanket FROST 0.
         assert g.get_workspace_size() == 0
-    elif route == "deselected":
-        assert g.selected_engine is None  # FROST barred -> the backend served it
-    ws = torch.empty(max(g.get_workspace_size(), 1), device="cuda", dtype=torch.uint8)
-    y = torch.empty(1, M, N, dtype=torch.bfloat16, device="cuda")
-    g.execute({A: a, B: b, bias: bias_t, Y: y}, ws)
-    torch.cuda.synchronize()
-    torch.testing.assert_close(y, ref, atol=1e-1, rtol=1e-2)
-
-
-@_GPU
-def test_build_walk_falls_through_a_declining_plan(caplog, monkeypatch):
-    """A plan that declines at build time is logged and the walk moves to the
-    next entry — here a python engine ranked ahead of the backend, so the graph
-    still builds and executes natively with no exception reaching the user
-    (a select_plan pin is strict instead; see build_plans)."""
-    import logging
-
-    from cudnn.engines import BaseEngine
-
-    class Boom(BaseEngine):
-        name = "frost_fake_always_fails"
-        engine_id = _FAKE + 1
-
-        def build_plan(self, graph, plan, ctx=None):
-            raise NotImplementedError("frost build boom")
-
-        def execute(self, graph, tensor_data, ctx=None):
-            raise AssertionError("should never run")
-
-    boom = Boom()
-    _offer(monkeypatch, boom)
-
-    a, b, bias_t, ref = _operands()
-    g, A, B, bias, Y = _build_matmul_bias_relu()
-    monkeypatch.setattr(heuristics, "rank", lambda graph, engines, backend_plans, modes=None: [PlanConfig(boom.engine_id)] + list(backend_plans))
-    _plan(g)
-    assert g.get_plan_name_at_index(0) == "frost_fake_always_fails"
-    g.check_support()
-    with caplog.at_level(logging.INFO, logger="cudnn.pygraph"):
-        g.build_plans()  # entry 0 declines -> the walk lands on the backend
-    assert any("declined at build time" in r.getMessage() for r in caplog.records)
-    assert g.selected_engine is None and is_backend_engine(g.plans[g._plan_index].engine_id)
     ws = torch.empty(max(g.get_workspace_size(), 1), device="cuda", dtype=torch.uint8)
     y = torch.empty(1, M, N, dtype=torch.bfloat16, device="cuda")
     g.execute({A: a, B: b, bias: bias_t, Y: y}, ws)
@@ -314,71 +269,6 @@ def test_frost_is_one_entry_of_the_ranked_list():
     assert sum(1 for p in g.plans if is_python_engine(p.engine_id)) == 1
 
 
-class _LoweredSpy:
-    """Records deselect_engines calls reaching the lowered C++ graph, forwarding
-    everything else untouched."""
-
-    def __init__(self, real):
-        self._real = real
-        self.deselected = []
-
-    def __getattr__(self, name):
-        attr = getattr(self._real, name)
-        if name != "deselect_engines":
-            return attr
-
-        def _record(names, *args, **kwargs):
-            self.deselected.append(list(names))
-            return attr(names, *args, **kwargs)
-
-        return _record
-
-
-def _spy_on_lowered(g):
-    spy = _LoweredSpy(g.__dict__["_lowered_graph"])
-    g.__dict__["_lowered_graph"] = spy
-    return spy
-
-
-def _native_engine_token(g):
-    """The cuDNN engine name (e.g. "eng0") of the leading NATIVE plan."""
-    return g.get_plan_name_at_index(_first_backend_index(g)).split("_")[0]
-
-
-@_GPU
-def test_deselect_native_engine_reaches_cudnn():
-    """deselect_engines is the classic API and stays a passthrough to the
-    lowered C++ graph — pygraph bars the name across the whole ranked list AND
-    forwards it, so a backend engine name still reaches cuDNN itself."""
-    g, *_ = _build_matmul_bias_relu()
-    _plan(g)
-    native = _native_engine_token(g)
-    spy = _spy_on_lowered(g)
-    assert g.deselect_engines([native]) is g  # fluent, like the other setters
-    assert spy.deselected == [[native]]
-
-
-@_GPU
-def test_deselect_mixed_frost_and_native():
-    """A mixed list bars the python engine locally AND forwards to cuDNN; the
-    barred FROST entry is skipped by the build walk."""
-    a, b, bias_t, ref = _operands()
-    g, A, B, bias, Y = _build_matmul_bias_relu()
-    _plan(g)
-    native = _native_engine_token(g)
-    spy = _spy_on_lowered(g)
-    g.deselect_engines([_FROST, native])
-    assert _FROST in g._barred_names
-    assert spy.deselected == [[_FROST, native]]
-    g.build_plans()
-    assert g.selected_engine is None  # FROST barred -> a backend plan serves it
-    ws = torch.empty(max(g.get_workspace_size(), 1), device="cuda", dtype=torch.uint8)
-    y = torch.empty(1, M, N, dtype=torch.bfloat16, device="cuda")
-    g.execute({A: a, B: b, bias: bias_t, Y: y}, ws)
-    torch.cuda.synchronize()
-    torch.testing.assert_close(y, ref, atol=1e-1, rtol=1e-2)
-
-
 def _build_moe(S=512, N=256, K=256, E=4):
     g = cudnn.pygraph(
         io_data_type=cudnn.data_type.BFLOAT16,
diff --git a/test/python/gemm/frost/test_matmul.py b/test/python/gemm/frost/test_matmul.py
index 9af675e33..71f5668f9 100644
--- a/test/python/gemm/frost/test_matmul.py
+++ b/test/python/gemm/frost/test_matmul.py
@@ -19,10 +19,14 @@
 import torch
 
 from gemm_test_utils import (
+    requires_int8_mma,
     requires_sm100,
     Plan as _plan,
     vp as _vp,
     resolve as _resolve,
+    e5m3_quant_ref as _e5m3_quant_ref,
+    e5m3_to_float as _e5m3_to_float,
+    requires_sm107,
 )
 
 # Module-wide GPU gate — every test here is end-to-end and needs a B200.
@@ -34,10 +38,6 @@
 from cudnn.gemm.frost.compiler import _current_arch, _epi_vec_bytes
 from cudnn.gemm.frost.tile_config import CATALOG, by_name
 
-# INT8 matmul runs only on SM 100 or SM 110 (disjoint range).
-_INT8_SM_RANGES = ((100, 101), (110, 111))
-
-
 _TORCH_DTYPE = {
     "bf16": torch.bfloat16,
     "fp16": torch.float16,
@@ -904,6 +904,12 @@ def _col_quant_reference(
     blocks = x.view(B, M // block_size, block_size, N)
     output_max = 448.0 if out_dtype is torch.float8_e4m3fn else 57344.0
     scale_f = blocks.abs().amax(dim=2) / output_max
+    if scale_dtype == "e5m3":
+        scale = _e5m3_quant_ref(scale_f)
+        back = _e5m3_to_float(scale.to(torch.int32))
+        inv = torch.where(back > 0, back.reciprocal(), 0.0)
+        q = (blocks * inv.unsqueeze(2)).clamp(-output_max, output_max)
+        return q.to(out_dtype).view(B, M, N), scale
     if scale_dtype is torch.float8_e8m0fnu:
         safe = torch.where(scale_f > 0, scale_f, 1.0)
         scale_f = torch.where(
@@ -978,6 +984,40 @@ def test_dense_col_block_scale_quant(config_name) -> None:
     torch.testing.assert_close(q.float(), q_ref.float(), atol=0, rtol=0)
 
 
+@requires_sm107
+@pytest.mark.parametrize("block_size", [16, 32])
+def test_dense_col_quant_e5m3_scale(block_size) -> None:
+    """COL block quantize with an E5M3 scale. This path keeps the scale in a
+    per-lane register (`_scale_mine_*`) before storing, so it exercises the
+    zero-init + assignment + store of the byte carrier that row quant does not."""
+    cfg, cta_group, scheduler = _resolve("CONFIG_sm100_128x128x128_128x128x32_cluster1x1_1ctamma")
+    M, N, K = 256, 128, 128
+    g = cudnn.pygraph(io_data_type=cudnn.data_type.BFLOAT16, intermediate_data_type=cudnn.data_type.FLOAT, compute_data_type=cudnn.data_type.FLOAT)
+    A = g.tensor(name="A", dim=[1, M, K], stride=_a_stride_batched(M, K, "k"))
+    B = g.tensor(name="B", dim=[1, K, N], stride=_b_stride_batched(N, K, "k"))
+    C = g.matmul(A=A, B=B, name="mm")
+    S = g.swish(input=C, name="sw")
+    Q, QS = g.block_scale_quantize(input=S, block_size=block_size, axis=1, name="q")
+    Q.set_output(True).set_data_type(cudnn.data_type.FP8_E4M3)
+    QS.set_dim([1, M // block_size, N]).set_stride([M // block_size * N, N, 1])
+    QS.set_output(True).set_data_type(cudnn.data_type.FP8_E5M3)
+
+    compiled = _plan(g, config=cfg, cta_group=cta_group, scheduler=scheduler)
+    assert compiled.chain.quants[0].axis == 1 and compiled.chain.quants[0].scale_dtype == "fp8_e5m3"
+
+    a, b, _ = _mkdata(M, N, K, "bf16", "bf16")
+    q = torch.empty(1, M, N, dtype=torch.float8_e4m3fn, device="cuda")
+    q_scale = torch.empty(1, M // block_size, N, dtype=torch.int8, device="cuda")
+    compiled(_vp(compiled, a, b, [q, q_scale]))
+    torch.cuda.synchronize()
+
+    ref_mm = torch.einsum("bmk,bnk->bmn", a.to(torch.float32), b.to(torch.float32))
+    ref_sw = ref_mm * torch.sigmoid(ref_mm)
+    q_ref, scale_ref = _col_quant_reference(ref_sw, block_size, torch.float8_e4m3fn, "e5m3")
+    torch.testing.assert_close(q_scale.view(torch.uint8).float(), scale_ref.float(), atol=0, rtol=0)
+    torch.testing.assert_close(q.float(), q_ref.float(), atol=0, rtol=0)
+
+
 def test_dense_row_col_dual_quant_f8_reorder() -> None:
     """The cutedsl dual-output pattern: one producer -> row quant + col quant,
     both with F8_128x4 scale reordering."""
@@ -2023,14 +2063,12 @@ def test_mixed_fp8_matmul(config_name: str, a_dt: str, b_dt: str) -> None:
 }
 
 
+@requires_int8_mma
 @pytest.mark.parametrize("config_name", _INT8_CONFIGS, ids=[_config_id(n) for n in _INT8_CONFIGS])
 @pytest.mark.parametrize("out_dt", list(_INT8_OUT_DTYPES))
 def test_int8_matmul(config_name: str, out_dt: str) -> None:
     """INT8×INT8→INT32, output ∈ {fp32,bf16,fp16,int32,fp8}; bit-exact vs a
     rounded integer reference (values small enough that the rounding is exact)."""
-    sm = _current_arch()
-    if sm is not None and not any(lo <= sm < hi for lo, hi in _INT8_SM_RANGES):
-        pytest.skip(f"int8 matmul unsupported on sm_{sm} (SM 100/110 only)")
     cfg, cta_group, scheduler = _resolve(config_name)
     M = N = K = 256
     cudnn_dt, torch_dt, vmax = _INT8_OUT_DTYPES[out_dt]
diff --git a/test/python/gemm/frost/test_matmul_mainloop_fusion.py b/test/python/gemm/frost/test_matmul_mainloop_fusion.py
index cff813a80..0a5522887 100644
--- a/test/python/gemm/frost/test_matmul_mainloop_fusion.py
+++ b/test/python/gemm/frost/test_matmul_mainloop_fusion.py
@@ -15,6 +15,7 @@
 import torch
 
 from gemm_test_utils import (
+    requires_int8_mma,
     requires_sm100,
     Plan as _plan,
     vp as _vp,
@@ -302,6 +303,7 @@ def test_e2e_mainloop_fp16(cfg) -> None:
     ],
 )
 @requires_sm100
+@requires_int8_mma
 def test_e2e_mainloop_int8(op, cfg) -> None:
     """INT8 mainloop fusion: f(int8 A) @ int8 B → int32 acc → fp32 out.
     Exercises the integer idesc + int32→fp32 widen; bit-exact vs int reference."""
diff --git a/test/python/gemm/frost/test_moe_grouped_block_scale_matmul_fwd.py b/test/python/gemm/frost/test_moe_grouped_block_scale_matmul_fwd.py
index c0a3538e6..fe3432e79 100644
--- a/test/python/gemm/frost/test_moe_grouped_block_scale_matmul_fwd.py
+++ b/test/python/gemm/frost/test_moe_grouped_block_scale_matmul_fwd.py
@@ -14,6 +14,7 @@
 
 from gemm_test_utils import (
     requires_sm100,
+    requires_sm107,
     Plan as _plan,
     vp_bs as _vp_bs,
     E2M1 as _E2M1,
@@ -27,6 +28,7 @@
     assert_block_scale_reduction_close as _assert_block_scale_reduction_close,
 )
 
+from cudnn.gemm.frost.dtypes import DTYPE_FROM_CUDNN as _DTYPE_FROM_CUDNN
 from cudnn.gemm.frost.compiler import jit_from_cudnn_graph
 from cudnn.gemm.frost.graph_analyzer import analyze
 from cudnn.gemm.frost.tile_config import by_name
@@ -244,7 +246,7 @@ def test_analyzer_detects_moe_grouped_block_scale_matmul_fwd() -> None:
     assert chain.has_moe and chain.has_block_scale
     assert chain.moe.num_experts == E
     assert chain.moe.mode == "none"
-    assert chain.block_scale.combo == "nvfp4"
+    assert (chain.block_scale.sf_dtype, chain.block_scale.block_size) == ("fp8_e4m3", 16)
     assert chain.matmul.a_dtype == "fp4_e2m1"
     assert chain.matmul.b_dtype == "fp4_e2m1"
     assert (chain.matmul.M, chain.matmul.N, chain.matmul.K) == (S, N, K)
@@ -361,7 +363,9 @@ def _run_e2e(
         config=cfg,
         cta_group=cta_group,
     )
-    assert compiled.chain.block_scale.combo == combo
+    _blk, _, _sf_dt = _COMBOS[combo]
+    _bs = compiled.chain.block_scale
+    assert (_bs.sf_dtype, _bs.block_size) == (_DTYPE_FROM_CUDNN[_sf_dt], _blk)
 
     # SFA reordered + padded to 128 rows PER GROUP, then concatenated (for
     # 128-aligned groups this equals a single global _to_blocked). SFB per-expert.
@@ -767,12 +771,81 @@ def test_auto_config_is_accepted_by_the_registry(S, N):
     """Same invariant as the dense block-scale case: the grouped path shares the
     BlockScaleSpec machinery, so its 128-multiple tile constraint applies too and
     ``select_config`` must not pick a geometry the registry rejects."""
-    from cudnn.gemm.frost.kernel_registry import candidates
-    from cudnn.gemm.frost.tile_config import select_config
+    from cudnn.gemm.frost.kernel_registry import candidates, preferred_pipeline
+    from cudnn.gemm.frost.tile_config import as_pipeline, select_config
 
     chain = analyze(_build_graph(8, S, N, 512, 8))
     assert chain.has_block_scale and chain.has_moe
     cfg, _cta_group, _sched = select_config(chain.matmul.M, chain.matmul.N, chain.num_gemms, block_scale=chain.has_block_scale)
+    cfg = as_pipeline(cfg, preferred_pipeline(chain))  # the config build_gemm_plan actually builds
     accepted = {c.name for _t, c in candidates(chain)}
     assert accepted, "the registry accepts no geometry at all for this chain"
     assert cfg.name in accepted, f"select_config picked {cfg.name!r}, which the registry rejects for this graph"
+
+
+# ---------------------------------------------------------------------------
+# sm107 pipeline (the sm100 grouped pipeline on the 64-byte-K block-scale MMA)
+# ---------------------------------------------------------------------------
+
+_SM107_CFG = "CONFIG_sm107_128x256x128_128x256x64_cluster2x1"
+_SM107_CFG_1CTA = "CONFIG_sm107_128x256x128_128x256x64_cluster1x1"
+
+
+def test_sm107_template_selection_and_arch_gate(monkeypatch) -> None:
+    from cudnn.gemm.frost import compiler as C
+    from cudnn.gemm.frost.kernel_registry import TEMPLATES, select_template
+
+    monkeypatch.setattr(C, "_current_arch", lambda: 107)
+    chain = analyze(_build_graph(2, 512, 256, 512, num_groups=2))
+    for cta_group, cfg_name, want in (
+        (1, _SM107_CFG_1CTA, "sm107_moe_grouped_block_scale_matmul_fwd_1ctamma.py"),
+        (2, _SM107_CFG, "sm107_moe_grouped_block_scale_matmul_fwd_2ctamma.py"),
+    ):
+        cfg = by_name(cfg_name)
+        tmpl = select_template(chain, cfg, cta_group=cta_group, scheduler="clc")
+        assert tmpl.file == want
+        assert tmpl.accepts(chain, cfg) is None
+    # An sm100 config still pairs with the sm100 grouped templates on the same GPU.
+    assert select_template(chain, by_name(_CFG), cta_group=2, scheduler="clc").file == "sm100_moe_grouped_block_scale_matmul_fwd_2ctamma.py"
+    # ... and the sm107 templates are gated off older Blackwell.
+    monkeypatch.setattr(C, "_current_arch", lambda: 100)
+    tmpl = next(t for t in TEMPLATES if t.file == "sm107_moe_grouped_block_scale_matmul_fwd_1ctamma.py")
+    assert "107 <= SM < 110" in tmpl.accepts(chain, by_name(_SM107_CFG_1CTA))
+
+
+@pytest.mark.parametrize("combo", ["nvfp4", "mxfp4", "mxfp8"])
+@pytest.mark.parametrize("cfg_name,cta_group", [(_SM107_CFG, 2), (_SM107_CFG_1CTA, 1)])
+@requires_sm107
+def test_e2e_sm107(combo, cfg_name, cta_group) -> None:
+    _run_e2e(
+        E=2,
+        S=1024,
+        N=256,
+        K=512,
+        offsets_list=[0, 256, 384, 512],
+        combo=combo,
+        config_name=cfg_name,
+        cta_group=cta_group,
+    )
+
+
+@pytest.mark.parametrize("combo", ["nvfp4", "mxfp8"])
+@pytest.mark.parametrize("cta_group", [1, 2])
+@pytest.mark.parametrize("cta_m,cta_n", [(128, 256), (256, 128), (256, 256)])
+@requires_sm107
+def test_e2e_sm107_multi_mma_m(combo, cta_group, cta_m, cta_n) -> None:
+    # The grouped pipeline with the CTA tile split along M. The per-group-padded
+    # SF blob is unchanged; what moves is the TMEM side, where SFA is indexed per
+    # M block and SFB is walked across N blocks (they only differ once a block
+    # count exceeds one, i.e. at cta_m/cta_n = 256).
+    cluster = "cluster1x1" if cta_group == 1 else "cluster2x1"
+    name = f"CONFIG_sm107_{cta_m}x{cta_n}x128_128x{cta_n}x64_{cluster}"
+    _run_e2e(E=4, S=512, N=256, K=256, offsets_list=[0, 128, 256, 384], combo=combo, config_name=name, cta_group=cta_group)
+
+
+@pytest.mark.parametrize("cfg_name,cta_group", [(_SM107_CFG, 2), (_SM107_CFG_1CTA, 1)])
+@requires_sm107
+def test_e2e_sm107_unaligned_groups(cfg_name, cta_group) -> None:
+    # Group offsets that are not 128-aligned — the per-group-padded SF blob
+    # layout is the sm100 one, so the 64-byte-K MMA must not disturb it.
+    _run_e2e(E=2, S=512, N=256, K=512, offsets_list=[0, 100, 300], config_name=cfg_name, cta_group=cta_group)
diff --git a/test/python/gemm/frost/test_moe_grouped_block_scale_matmul_fwd_swiglu.py b/test/python/gemm/frost/test_moe_grouped_block_scale_matmul_fwd_swiglu.py
index c79312db6..af5b01a46 100644
--- a/test/python/gemm/frost/test_moe_grouped_block_scale_matmul_fwd_swiglu.py
+++ b/test/python/gemm/frost/test_moe_grouped_block_scale_matmul_fwd_swiglu.py
@@ -18,6 +18,7 @@
 
 from gemm_test_utils import (
     requires_sm100,
+    requires_sm107,
     Plan as _plan,
     E2M1 as _E2M1,
     to_blocked as _to_blocked,
@@ -26,6 +27,7 @@
     block_quant_ref as _block_quant_ref,
 )
 
+from cudnn.gemm.frost.dtypes import DTYPE_FROM_CUDNN as _DTYPE_FROM_CUDNN
 from cudnn.gemm.frost.graph_analyzer import analyze
 from cudnn.gemm.frost.tile_config import by_name
 
@@ -57,10 +59,13 @@ def _vp_moe_bs_mg(compiled, gemm_pairs, fto, outs, *aux):
 
 
 # cta_tile_n=128: dual block-scale TMEM fits two accs + SF only at n<=128.
-# (config, cta_group): 2-CTA cluster2x1 (reference) + 1-CTA cluster1x1.
+# (config, cta_group): 2-CTA cluster2x1 (reference) + 1-CTA cluster1x1, on both
+# the sm100 pipeline and the sm107 one (same geometry, 64-byte-K MMA).
 _GEOMETRIES = [
     ("CONFIG_sm100_128x128x128_128x128x32_cluster2x1", 2),
     ("CONFIG_sm100_128x128x128_128x128x32_cluster1x1", 1),
+    pytest.param("CONFIG_sm107_128x128x128_128x128x64_cluster2x1", 2, marks=requires_sm107),
+    pytest.param("CONFIG_sm107_128x128x128_128x128x64_cluster1x1", 1, marks=requires_sm107),
 ]
 
 
@@ -173,7 +178,7 @@ def test_analyzer_detects_dual_moe_grouped_block_scale_matmul_fwd() -> None:
     assert chain.has_moe and chain.has_block_scale and chain.is_multi_gemm
     assert chain.num_gemms == 2
     assert chain.num_a_operands == 1 and chain.num_b_operands == 2
-    assert chain.block_scale.combo == "nvfp4"
+    assert (chain.block_scale.sf_dtype, chain.block_scale.block_size) == ("fp8_e4m3", 16)
     assert chain.moe.num_experts == 2
     assert [o.op for o in chain.ops] == ["swish", "mul", "mul"]
     assert len(chain.outputs) == 1 and chain.outputs[0].source == "op_2"
@@ -255,7 +260,9 @@ def test_dual_moe_grouped_block_scale_matmul_fwd_swiglu(combo, cfg_name, cta_gro
 
     cfg = by_name(cfg_name)
     compiled = _plan(_build_graph(E, S, N, K, num_groups, combo), config=cfg, cta_group=cta_group)
-    assert compiled.chain.block_scale.combo == combo
+    _blk, _, _sf_dt = _COMBOS[combo]
+    _bs = compiled.chain.block_scale
+    assert (_bs.sf_dtype, _bs.block_size) == (_DTYPE_FROM_CUDNN[_sf_dt], _blk)
 
     # SFA padded to 128 rows PER GROUP, then concatenated; SFB per-expert.
     sfa_parts = [_to_blocked(sfa_log[offsets_list[gi] : offsets_list[gi + 1] if gi + 1 < num_groups else S]) for gi in range(num_groups)]
diff --git a/test/python/gemm/frost/test_tile_select_analytic.py b/test/python/gemm/frost/test_tile_select_analytic.py
index 2ef2dea8d..6d0732133 100644
--- a/test/python/gemm/frost/test_tile_select_analytic.py
+++ b/test/python/gemm/frost/test_tile_select_analytic.py
@@ -99,3 +99,23 @@ def test_moe_and_mainloop_never_get_static():
         for N in (1024, 8192):
             _, _, sched = select_config(M, N, 1, K=4096, supports_static=False)
             assert sched == "clc"
+
+
+def test_a_new_pipeline_must_register_its_hardware_facts():
+    """A family that registers a config class but forgets a per-pipeline table
+    must raise, not inherit another family's value: the tables are hardware
+    facts, and a wrong MMA-inst K renders a descriptor that is silently wrong."""
+    import dataclasses
+
+    from cudnn.gemm.frost import tile_config as tc
+
+    @dataclasses.dataclass(frozen=True)
+    class ConfigSmFake(tc.TileConfig):
+        pass
+
+    tc._CONFIG_CLASS_BY_PIPELINE["sm_fake"] = ConfigSmFake
+    try:
+        with pytest.raises(NotImplementedError, match="MMA-inst K width not known for pipeline"):
+            tc.as_pipeline(tc.DEFAULT_CONFIG, "sm_fake")
+    finally:
+        del tc._CONFIG_CLASS_BY_PIPELINE["sm_fake"]
diff --git a/test/python/test_dispatch.py b/test/python/test_dispatch.py
index 1a69aa0a2..64e07ffaf 100644
--- a/test/python/test_dispatch.py
+++ b/test/python/test_dispatch.py
@@ -459,9 +459,7 @@ def test_mixed_ranking_dispatch(monkeypatch):
 def test_pinned_plan_that_declines_raises(monkeypatch):
     """A select_plan() pin is STRICT: the walk starts there and a decline raises
     instead of quietly running a different plan. Without a pin the same decline
-    only advances the walk (the GPU counterpart is
-    test_build_walk_falls_through_a_declining_plan in
-    test/python/gemm/frost/test_frontend_integration.py)."""
+    only advances the walk."""
 
     class Declines(BaseEngine):
         name = "declines_at_build"
@@ -481,6 +479,39 @@ def execute(self, graph, tensor_data, ctx=None):
         g.build_plans()
 
 
+def test_unpinned_plan_that_declines_advances_the_walk(monkeypatch, caplog):
+    """The other half of the pin rule: WITHOUT a pin the same decline is logged
+    and the walk moves to the next entry, so the graph still builds and no
+    exception reaches the user."""
+    import logging
+
+    from cudnn.engines import PlanConfig
+
+    class Declines(BaseEngine):
+        name = "declines_at_build"
+        engine_id = _FAKE + 81
+
+        def build_plan(self, graph, plan, ctx=None):
+            raise NotImplementedError("cannot compile this graph")
+
+        def execute(self, graph, tensor_data, ctx=None):
+            raise AssertionError("should never run")
+
+    declines, works = Declines(), _mk_engine(82)
+    _offer(monkeypatch, declines, works)
+    _ranking(monkeypatch, lambda graph, engines, backend_plans, modes=None: [PlanConfig(declines.engine_id), PlanConfig(works.engine_id)])
+
+    g = pygraph()
+    g.matmul(torch.randn(2, 2), torch.randn(2, 2))
+    g.create_execution_plans()
+    assert _plan_names(g) == ["declines_at_build", works.name]
+    with caplog.at_level(logging.INFO, logger="cudnn.pygraph"):
+        g.build_plans()
+    assert any("declined at build time" in r.getMessage() for r in caplog.records)
+    assert g.selected_engine is works
+    assert _plan_names(g)[g._plan_index] == works.name
+
+
 def test_empty_ranking_output_rejected(monkeypatch):
     """Ranking that returns [] is an error — there is no legal empty planning
     state (it would defeat the one-shot flag and every needs-planning check)."""