[#15853][fix] Gate B12x MoE on CUDA13 CuTe DSL

tensorrt_llm/_torch/moe/fused_moe/MOE_DEVELOPER_GUIDE.md

@@ -692,6 +692,14 @@ tensor cores. This is what makes Cutlass the landing spot when a W4A16 NVFP4
 layer finds no specialized backend — the B12x W4A16 leaf without flashinfer,
 or any SM outside Marlin's 89-99 and B12x's 120/121.
 
+On SM120/121 the B12x NVFP4 leaf `CuteDslB12xNvfp4FusedMoE` additionally needs
+`MoEDep.CUTEDSL_CUDA13`: the active `nvidia-cutlass-dsl` native payload must
+report CUDA 13 or newer through `cutlass.CUDA_VERSION`. The CUDA 12.x payload
+lowers the SM12x NVFP4 MMA atom to PTX that ptxas rejects, so without it the
+leaf is `DEP_MISSING` and falls back to Cutlass instead of failing in the
+decode-kernel JIT. The W4A16 NVFP4 leaf does not instantiate that atom and is
+not gated.
+
 ‡ No CuteDSL leaf publishes `fp8_block_scales`; the DeepGEMM and TRTLLM-Gen
 leaves own the format on SM100/103/107 and Cutlass on SM90/SM120. Because a
 `CUTEDSL` request considers only the CuteDSL family plus the Cutlass fallback,

tensorrt_llm/_torch/moe/fused_moe/fused_moe_cute_dsl_b12x.py

@@ -253,6 +253,20 @@ def can_implement(cls, p: MoEProblem, d: MoEDeployment) -> MoEEligibility:
                 MoERejectReason.DEP_MISSING,
                 f"{cls.__name__} requires the flashinfer package",
             )
+        # CUDA 12.x CuTe DSL lowers the SM12x NVFP4 MMA atom to the internal
+        # ``_mma.block_scale...`` spelling, which ptxas rejects. CUDA 13.x
+        # emits the public ``mma.sync.aligned...kind::mxf4nvf4`` opcode. Only
+        # the pure-NVFP4 decode kernel instantiates that atom.
+        if cls.b12x_quant_algo == QuantAlgo.NVFP4 and not d.env.has_dep(MoEDep.CUTEDSL_CUDA13):
+            return _reject(
+                MoERejectReason.DEP_MISSING,
+                f"{cls.__name__} requires the active nvidia-cutlass-dsl "
+                "CUDA 13 native payload. CUDA 12.x CuTe DSL lowers FlashInfer's "
+                "B12x NVFP4 MMA to PTX that ptxas rejects with Unexpected "
+                "instruction types specified for '_mma'. Install "
+                "nvidia-cutlass-dsl[cu13] so cutlass.CUDA_VERSION reports CUDA "
+                "13.x or newer.",
+            )
         if d.eplb_enabled:
             return _reject(
                 MoERejectReason.EPLB_UNSUPPORTED,

tensorrt_llm/_torch/moe/fused_moe/impl_environment.py

@@ -39,6 +39,11 @@ class MoEDep(str, Enum):
     DEEPGEMM_MEGAMOE = "deepgemm_megamoe"
     #: ``nvidia-cutlass-dsl[cu13]`` is new enough for the MegaMoE CuteDSL ABI.
     MEGAMOE_CUTEDSL_RUNTIME = "megamoe_cutedsl_runtime"
+    #: The active ``nvidia-cutlass-dsl`` native payload targets CUDA 13 or
+    #: newer. Gates the NVFP4 decode kernel of ``CuteDslB12xFusedMoE``: the
+    #: CUDA 12.x payload lowers the SM12x NVFP4 MMA atom to PTX that ptxas
+    #: rejects.
+    CUTEDSL_CUDA13 = "cutedsl_cuda13"
     #: The ``trtllm::cute_dsl_megamoe_nvfp4_*`` custom ops are registered.
     MEGAMOE_CUTEDSL_OP = "megamoe_cutedsl_op"
     #: The installed CuTe DSL exposes the Rubin helpers the SM107 CuteDSL
@@ -62,6 +67,10 @@ class MoEEnvFlag(str, Enum):
     SHARED_EXPERT_FUSION = "TLLM_MOE_ENABLE_SHARED_EXPERT_FUSION"
 
 
+# Oldest CUDA major whose CuTe DSL payload emits the public
+# ``mma.sync.aligned...kind::mxf4nvf4`` opcode for the SM12x NVFP4 MMA atom.
+_MIN_CUTEDSL_CUDA_MAJOR = 13
+
 # Probe details are logged but excluded from the stable fingerprint.
 DepProbe = Callable[[], Tuple[bool, str]]
 
@@ -106,6 +115,24 @@ def _probe_megamoe_cutedsl_runtime() -> Tuple[bool, str]:
     return bool(available), "" if available else str(reason)
 
 
+def _probe_cutedsl_cuda13() -> Tuple[bool, str]:
+    try:
+        import cutlass
+    except ImportError as exc:
+        return False, f"import cutlass failed: {exc}"
+    # ``cutlass.CUDA_VERSION`` is the public spelling of
+    # ``cutlass.base_dsl.version_info.CUDA_VERSION``. It reports the CUDA
+    # version of the native payload that is actually loaded, which can still
+    # be 12.x with ``nvidia-cutlass-dsl-libs-cu13`` installed when the base
+    # payload shadows it.
+    major = getattr(getattr(cutlass, "CUDA_VERSION", None), "major", None)
+    if not isinstance(major, int):
+        return False, "cutlass.CUDA_VERSION.major is unavailable"
+    if major < _MIN_CUTEDSL_CUDA_MAJOR:
+        return False, f"active CuTe DSL payload targets CUDA {major}.x"
+    return True, ""
+
+
 def _probe_cutedsl_rubin() -> Tuple[bool, str]:
     from ...cute_dsl_utils import IS_CUTLASS_DSL_RUBIN_AVAILABLE
 
@@ -137,6 +164,7 @@ def _probe_megamoe_cutedsl_op() -> Tuple[bool, str]:
     MoEDep.DEEPGEMM_MEGAMOE: _probe_deepgemm_megamoe,
     MoEDep.MEGAMOE_CUTEDSL_RUNTIME: _probe_megamoe_cutedsl_runtime,
     MoEDep.MEGAMOE_CUTEDSL_OP: _probe_megamoe_cutedsl_op,
+    MoEDep.CUTEDSL_CUDA13: _probe_cutedsl_cuda13,
     MoEDep.CUTEDSL_RUBIN: _probe_cutedsl_rubin,
     MoEDep.LOCALITY_DOMAIN: _probe_locality_domain,
 }

tests/unittest/_torch/moe/test_cute_dsl_b12x_moe_backend.py

@@ -22,6 +22,7 @@
 import pytest
 import torch
 
+from tensorrt_llm._torch.model_config import ModelConfig
 from tensorrt_llm._torch.moe.fused_moe.fused_moe_cute_dsl_b12x import (
     CuteDslB12xFusedMoE,
     CuteDslB12xNvfp4FusedMoE,
@@ -35,7 +36,12 @@
     MoERejectReason,
     canonical_quant,
 )
-from tensorrt_llm._torch.moe.fused_moe.impl_environment import MoEDep
+from tensorrt_llm._torch.moe.fused_moe.impl_environment import MoEDep, _probe_cutedsl_cuda13
+from tensorrt_llm._torch.moe.fused_moe.moe_resolution import (
+    backend_family_of,
+    impl_class_for,
+    resolve_moe_impl,
+)
 from tensorrt_llm._torch.moe.fused_moe.quantization import (
     NVFP4CuteDslB12xFusedMoEMethod,
     NVFP4CutlassFusedMoEMethod,
@@ -65,13 +71,19 @@ def _deployment(
     sm: int,
     *,
     flashinfer: bool = True,
+    cutedsl_cuda13: bool = True,
     ep_size: int = 1,
     tp_size: int = 1,
     use_dp: bool = False,
     parallel_size: Optional[int] = None,
     eplb: bool = False,
 ) -> MoEDeployment:
     """Declare the machine rather than patching the probes that read it."""
+    deps = []
+    if flashinfer:
+        deps.append(MoEDep.FLASHINFER.value)
+    if cutedsl_cuda13:
+        deps.append(MoEDep.CUTEDSL_CUDA13.value)
     return MoEDeployment(
         ep_size=ep_size,
         tp_size=tp_size,
@@ -81,7 +93,7 @@ def _deployment(
         eplb_enabled=eplb,
         env=MoEEnvironment(
             sm=sm,
-            available_deps=(MoEDep.FLASHINFER.value,) if flashinfer else (),
+            available_deps=tuple(deps),
         ),
     )
 
@@ -174,6 +186,91 @@ def test_can_implement_rejects_missing_flashinfer():
     assert verdict.reject_reason is MoERejectReason.DEP_MISSING
 
 
+@pytest.mark.parametrize("sm_version", SUPPORTED_SM)
+def test_can_implement_rejects_nvfp4_without_cuda13_cute_dsl(sm_version):
+    """CUDA 12.x CuTe DSL emits PTX that ptxas rejects for the NVFP4 MMA."""
+    verdict = NVFP4_LEAF.can_implement(
+        _problem(), _deployment(sm_version, cutedsl_cuda13=False)
+    )
+    assert not verdict.eligible
+    assert verdict.reject_reason is MoERejectReason.DEP_MISSING
+    assert "CUDA 13 native payload" in verdict.detail
+    assert "Unexpected instruction types" in verdict.detail
+
+
+@pytest.mark.parametrize("sm_version", SUPPORTED_SM)
+def test_can_implement_cuda13_gate_is_limited_to_pure_nvfp4(sm_version):
+    """Only the pure-NVFP4 decode kernel instantiates the NVFP4 MMA atom."""
+    verdict = CuteDslB12xW4a16Nvfp4FusedMoE.can_implement(
+        _problem(QuantAlgo.W4A16_NVFP4), _deployment(sm_version, cutedsl_cuda13=False)
+    )
+    assert verdict.eligible
+
+
+def _cutedsl_model_config() -> ModelConfig:
+    cfg = ModelConfig()
+    cfg.moe_backend = "CUTEDSL"
+    cfg.quant_config = QuantConfig(quant_algo=QuantAlgo.NVFP4)
+    return cfg
+
+
+@pytest.mark.parametrize("sm_version", SUPPORTED_SM)
+def test_cutedsl_degrades_to_cutlass_without_cuda13_cute_dsl(sm_version):
+    """CUTEDSL + NVFP4 + SM12x avoids the unsupported CuTe DSL JIT path."""
+    report = resolve_moe_impl(
+        _cutedsl_model_config(),
+        problem=_problem(),
+        deployment=_deployment(sm_version, cutedsl_cuda13=False),
+    )
+    assert backend_family_of(impl_class_for(report)) == "CUTLASS"
+    assert report.degraded
+    b12x_rejections = [r for r in report.rejected if r.legacy_backend == NVFP4_LEAF.__name__]
+    assert len(b12x_rejections) == 1
+    assert b12x_rejections[0].reason is MoERejectReason.DEP_MISSING
+    assert "CUDA 13 native payload" in b12x_rejections[0].detail
+
+
+@pytest.mark.parametrize("sm_version", SUPPORTED_SM)
+def test_cutedsl_selects_b12x_with_cuda13_cute_dsl(sm_version):
+    report = resolve_moe_impl(
+        _cutedsl_model_config(),
+        problem=_problem(),
+        deployment=_deployment(sm_version),
+    )
+    assert impl_class_for(report) is NVFP4_LEAF
+    assert not report.degraded
+
+
+def test_probe_cutedsl_cuda13_reports_missing_cutlass(monkeypatch):
+    # A ``None`` entry in ``sys.modules`` makes ``import cutlass`` raise
+    # ImportError without touching an installed package.
+    monkeypatch.setitem(sys.modules, "cutlass", None)
+    available, detail = _probe_cutedsl_cuda13()
+    assert not available
+    assert "import cutlass failed" in detail
+
+
+@pytest.mark.parametrize(
+    ("cuda_version", "expected"),
+    [
+        (None, False),
+        (types.SimpleNamespace(), False),
+        (types.SimpleNamespace(major="13"), False),
+        (types.SimpleNamespace(major=12), False),
+        (types.SimpleNamespace(major=13), True),
+        (types.SimpleNamespace(major=14), True),
+    ],
+)
+def test_probe_cutedsl_cuda13_reads_public_cutlass_api(monkeypatch, cuda_version, expected):
+    cutlass_module = types.ModuleType("cutlass")
+    if cuda_version is not None:
+        cutlass_module.CUDA_VERSION = cuda_version
+    monkeypatch.setitem(sys.modules, "cutlass", cutlass_module)
+    available, detail = _probe_cutedsl_cuda13()
+    assert available is expected
+    assert bool(detail) is not expected
+
+
 def test_can_implement_rejects_eplb():
     """EPLB is its own reject class, not a topology one: the machine is fine."""
     verdict = NVFP4_LEAF.can_implement(_problem(), _deployment(120, eplb=True))

tests/unittest/_torch/moe/test_moe_impl.py

@@ -1981,7 +1981,9 @@ def test_cutlass_build_gate_does_not_answer_for_an_unsupported_sm():
                 use_dp=False,
                 num_slots=8,
                 env=MoEEnvironment(
-                    sm=120, oss_cutlass_moe=False, available_deps=(MoEDep.FLASHINFER,)
+                    sm=120,
+                    oss_cutlass_moe=False,
+                    available_deps=(MoEDep.FLASHINFER, MoEDep.CUTEDSL_CUDA13),
                 ),
             ),
             MoERejectReason.TOPOLOGY_UNSUPPORTED,
@@ -2653,7 +2655,11 @@ def test_b12x_cutlass_build_mode_resolution(is_oss, sm, quant_algo, activation,
             use_dp=False,
             num_slots=8,
             force_dynamic_quantization=dynamic,
-            env=MoEEnvironment(sm=sm, oss_cutlass_moe=is_oss, available_deps=(MoEDep.FLASHINFER,)),
+            env=MoEEnvironment(
+                sm=sm,
+                oss_cutlass_moe=is_oss,
+                available_deps=(MoEDep.FLASHINFER, MoEDep.CUTEDSL_CUDA13),
+            ),
         ),
     )
     rejection = not is_oss and quant_algo == QuantAlgo.NVFP4 and (activation == "Relu2" or dynamic)