`nvidia-cutlass-dsl` 4.5.0: `nvvm.mma.block_scale` lowering produces PTX rejected by ptxas (sm_120/120f/121a)

[idonati]

Summary

In nvidia-cutlass-dsl==4.5.0 (the stable release that fixed the MmaSM120BlockScaledOp admissibility check for sm_121a), the MLIR-to-PTX lowering of nvvm.mma.block_scale for FP4 (E2M1) inputs with UE4M3 scales produces PTX that CUDA 13.0 ptxas rejects with:

ptxas application ptx input, line 969; error : Unexpected instruction types specified for '_mma'
...50+ instructions, same error...

The hand-written equivalent mma.sync.aligned.m16n8k64.row.col.kind::mxf4nvf4.block_scale.scale_vec::4X.f32.e2m1.e2m1.f32.ue4m3 compiles cleanly on .target sm_120a, .target sm_120f, and .target sm_121a. So the lowering pipeline is emitting something different — likely a slightly-different LLVM intrinsic encoding — that the user-facing ptxas can't decode.

Environment

• nvidia-cutlass-dsl==4.5.0 (also nvidia-cutlass-dsl-libs-base==4.5.0, nvidia-cutlass-dsl-libs-cu13==4.5.0)
• CUDA 13.0 toolkit (ptxas built Wed_Aug_20_01:53:56_PM_PDT_2025)
• Hardware: NVIDIA GB10 (sm_121a) — DGX Spark
• Triggered via flashinfer-python==0.6.11 (the b12x_fused_moe backend's NVFP4 path)
• Also reproduced via vLLM PR Integrate flashinfer b12x MoE and FP4 GEMM kernels for SM120/121 vllm-project/vllm#40082, --moe-backend flashinfer_b12x

Reproduction

The full IR dump is large (~2000 lines) because ptxas flags ~50 mma instructions per kernel — happy to attach if helpful. The offending MLIR op is:

%X = "nvvm.mma.block_scale"(%a, %sa_data, %sa_byteid_lo, %sa_byteid_hi,
                             %b, %sb_data, %sb_byteid_lo, %sb_byteid_hi, %c) <{
    aType = #nvvm.mma_type<e2m1>,
    bType = #nvvm.mma_type<e2m1>,
    blockScaleFormat = #nvvm.block_scale_format<ue4m3>,
    cType = #nvvm.mma_type<f32>,
    layoutA = #nvvm.mma_layout<row>,
    layoutB = #nvvm.mma_layout<col>,
    scaleVecSize = #nvvm.scale_vec_size<x4>,
    shape = #nvvm.shape<m = 16, n = 8, k = 64>
}> : (vector<4xi32>, i32, i16, i16, vector<2xi32>, i32, i16, i16, vector<4xf32>) -> vector<4xf32>

With target attribute:

"gpu.module"() <{sym_name = "kernels", targets = [#nvvm.target<chip = "sm_121a", flags = {"ptx-cmd-options" = []}>]}> ({...})

Switching the target to sm_120a, sm_120f, or sm_121a produces the same ptxas error.

Hand-written PTX (compiles fine on all three targets)

.version 8.8
.target sm_121a       // or sm_120a or sm_120f — all work
.address_size 64
.visible .entry test() {
  .reg .b32 a<4>, b<2>, c<4>, d<4>;
  .reg .b32 sa, sb;
  mov.b32 a0, 0; mov.b32 a1, 0; mov.b32 a2, 0; mov.b32 a3, 0;
  mov.b32 b0, 0; mov.b32 b1, 0;
  mov.b32 c0, 0; mov.b32 c1, 0; mov.b32 c2, 0; mov.b32 c3, 0;
  mov.b32 sa, 0; mov.b32 sb, 0;
  mma.sync.aligned.m16n8k64.row.col.kind::mxf4nvf4.block_scale.scale_vec::4X.f32.e2m1.e2m1.f32.ue4m3
    {d0,d1,d2,d3}, {a0,a1,a2,a3}, {b0,b1}, {c0,c1,c2,c3}, sa, {0,0}, sb, {0,0};
  ret;
}

$ ptxas test.ptx -o test.cubin --gpu-name=sm_121a; echo exit=$?
exit=0

So ptxas accepts this exact instruction at every Blackwell consumer-card target. The cute-dsl emission must be producing slightly different operand-encoding, register-encoding, or intrinsic name that ptxas can't decode (note ptxas calls it _mma, with underscore — possibly the LLVM-intrinsic name leaking through rather than the user-facing mma.sync mnemonic).

Smaller related bug

base_dsl/runtime/cuda.py:_get_gpu_arch_info has no (12, 1) entry in gpu_arch_map:

gpu_arch_map = {
    ...
    (12, 0): ("Blackwell", "sm_120a", ["sm_120a"]),  # RTX PRO 6000 / RTX 50 Series
    # missing: (12, 1) for GB10 / DGX Spark
}
return gpu_arch_map.get((major, minor),
                        ("Unknown", f"sm_{major}{minor}", [f"sm_{major}{minor}"]))

GB10 falls through to ("Unknown", "sm_121", ["sm_121"]) — note sm_121 without the a suffix. Users have to manually set CUTE_DSL_ARCH=sm_121a to override.

Suggested one-line fix:

(12, 1): ("Blackwell", "sm_121a", ["sm_121a", "sm_120f"]),  # GB10 / DGX Spark

Why this matters

Multiple downstream issues are blocked on this:

• Integrate flashinfer b12x MoE and FP4 GEMM kernels for SM120/121 vllm-project/vllm#40082 (b12x MoE / FP4 GEMM kernels for sm_120/121) — cannot land usefully until cute-dsl emits valid PTX for this MMA kind
• [Bug]: Xiaomi MiMo v2.5 broken on SM12x vllm-project/vllm#41519 (shadowlilac-oss's mimo_v2 repro on RTX PRO 6000 sm_120) — same root cause downstream
• [Bug] NVFP4 MoE models crash on GB10 (SM121) during CUDA graph capture flashinfer-ai/flashinfer#2776 (NVFP4 MoE crash on sm_120/121)
• NVIDIA forum thread "MiMo-V2.5 (New model)" (multiple consumer-Blackwell users)

Note that the precompiled-cubin alternative path (--moe-backend flashinfer_cutlass in vLLM) is ALSO broken on consumer Blackwell, but for a different reason: flashinfer-cubin==0.6.11 ships 12,681 FP4 cubins targeting Sm100a/Sm100f/Sm103a only — zero sm_120 or sm_121 cubins. So neither the JIT path (this bug) nor the AOT path (missing wheel coverage) currently works. Closing this one would unblock the JIT path, at minimum.

Happy to provide full reproducer (Dockerfile + recipe + IR dump) if useful for triage.

[depaulmillz]

When you updated to 4.5.0, did you fully uninstall the prior wheel? There is a known issue where this can lead to issues compiling like above.

[idonati]

Thanks for the fast response @depaulmillz — really appreciated.

Tried it. Same error.

Built a fresh image with an explicit clean uninstall before installing 4.5.0 from re-downloaded wheels (no --force-reinstall, nothing left behind from any prior wheel). Dockerfile:

FROM <base>
RUN pip uninstall -y nvidia-cutlass-dsl nvidia-cutlass-dsl-libs-base nvidia-cutlass-dsl-libs-cu13 \
 && rm -rf /usr/local/lib/python3.12/dist-packages/nvidia_cutlass_dsl* \
 && rm -rf /usr/local/lib/python3.12/dist-packages/cutlass
COPY wheels/ /tmp/wheels/
RUN pip install --no-deps \
      /tmp/wheels/nvidia_cutlass_dsl_libs_base-4.5.0-cp312-cp312-manylinux_2_28_aarch64.whl \
      /tmp/wheels/nvidia_cutlass_dsl_libs_cu13-4.5.0-cp312-cp312-manylinux_2_28_aarch64.whl \
      /tmp/wheels/nvidia_cutlass_dsl-4.5.0-py3-none-any.whl

Verified inside the running container:

$ ls -d /usr/local/lib/python3.12/dist-packages/nvidia_cutlass_dsl*
/usr/local/lib/python3.12/dist-packages/nvidia_cutlass_dsl
/usr/local/lib/python3.12/dist-packages/nvidia_cutlass_dsl-4.5.0.dist-info
/usr/local/lib/python3.12/dist-packages/nvidia_cutlass_dsl.pth
/usr/local/lib/python3.12/dist-packages/nvidia_cutlass_dsl_libs_base-4.5.0.dist-info
/usr/local/lib/python3.12/dist-packages/nvidia_cutlass_dsl_libs_cu13-4.5.0.dist-info

$ pip show nvidia-cutlass-dsl nvidia-cutlass-dsl-libs-base nvidia-cutlass-dsl-libs-cu13 | grep -E "^(Name|Version)"
Name: nvidia-cutlass-dsl
Version: 4.5.0
Name: nvidia-cutlass-dsl-libs-base
Version: 4.5.0
Name: nvidia-cutlass-dsl-libs-cu13
Version: 4.5.0

Re-launched the same workload on 8× DGX Spark (sm_121a). All 8 ranks loaded weights, FlashInfer dispatched FLASHINFER_B12X, and at the first MoE JIT we hit the same error at the same PTX line:

File ".../nvidia_cutlass_dsl/python_packages/cutlass/base_dsl/dsl.py", line 1385, in compile_and_jit
File ".../nvidia_cutlass_dsl/python_packages/cutlass/base_dsl/compiler.py", line 145, in compile
  pm.run(module.operation)
cutlass._mlir._mlir_libs._site_initialize.<locals>.MLIRError: Failure while executing pass pipeline:
error: unknown: NVPTX compiler invocation failed, error log:
  ptxas application ptx input, line 969; error : Unexpected instruction types specified for '_mma'
  ptxas application ptx input, line 970; error : Unexpected instruction types specified for '_mma'
  ...
  ptxas application ptx input, line 1111; error : Unexpected instruction types specified for '_mma'

Same module attribute targets = [#nvvm.target<chip = "sm_121a", flags = {"ptx-cmd-options" = []}>] as before.

So it doesn't look like wheel-layering cruft for us — appears to be a real lowering issue. Wheel SHA256s for reference:

3b051fe02ca69422ab840e64d9865667aba288a3984a7ca4ccd038a82aef1344  nvidia_cutlass_dsl-4.5.0-py3-none-any.whl
3f7c133d31fa82ae7db697fd6943a5f9a2c97c8a40ee1056c67ef29fe00974d8  nvidia_cutlass_dsl_libs_base-4.5.0-cp312-cp312-manylinux_2_28_aarch64.whl
a799eb9d65ba03095444907b8bf617b5a8bd7d03d1b95cec9637c558af6d0b60  nvidia_cutlass_dsl_libs_cu13-4.5.0-cp312-cp312-manylinux_2_28_aarch64.whl

(All three downloaded fresh from PyPI in the build — happy to also try pypi.nvidia.com if those differ.)

Happy to upload the full IR dump (the dumped MLIR module at the point of failure is ~2000 lines) or run any targeted diagnostic you'd find useful. Reproducer is the b12x backend in vllm-project/vllm#40082 on consumer Blackwell — but we're seeing the same error on plain chip = "sm_120f" family target too, so it doesn't appear sm_121a-specific.

[depaulmillz]

Thanks for reporting. I have been able to reproduce on a Spark. The CUDA 13 wheel works, but it appears the 12.9 wheel fails during compilation.

[idonati]

Thanks for the repro! That's encouraging — confirms it's not just our cluster.

The CUDA 13 wheel works, but it appears the 12.9 wheel fails during compilation.

I should have led with this: we are on the cu13 wheel. Full snapshot from inside the failing container:

$ python3 -c 'import sys, platform; print(sys.version); print(platform.platform(), platform.machine())'
3.12.3 (main, Aug 14 2025, 17:47:21) [GCC 13.3.0]
Linux-6.17.0-1014-nvidia-aarch64-with-glibc2.39 aarch64

$ pip list | grep -iE '^nvidia-|^cuda|^torch|^cutlass'
cuda-bindings                            13.2.0
cuda-pathfinder                          1.5.4
cuda-python                              13.2.0
cuda-tile                                1.3.0
cuda-toolkit                             13.0.2
nvidia-cublas                            13.1.0.3
nvidia-cuda-cupti                        13.0.85
nvidia-cuda-nvrtc                        13.0.88
nvidia-cuda-runtime                      13.0.96
nvidia-cudnn-cu13                        9.19.0.56
nvidia-cudnn-frontend                    1.15.0
nvidia-cufft                             12.0.0.61
nvidia-cufile                            1.15.1.6
nvidia-curand                            10.4.0.35
nvidia-cusolver                          12.0.4.66
nvidia-cusparse                          12.6.3.3
nvidia-cusparselt-cu13                   0.8.0
nvidia-cutlass-dsl                       4.5.0
nvidia-cutlass-dsl-libs-base             4.5.0
nvidia-cutlass-dsl-libs-cu13             4.5.0
nvidia-nccl-cu13                         2.30.4
nvidia-nvjitlink                         13.0.88
nvidia-nvtx                              13.0.85
torch                                    2.11.0

Everything is cu13. There is no nvidia-cutlass-dsl-libs-cu12 package on PyPI (only cu13 exists at 4.5.0), so we're not accidentally on a cu12 build. ptxas is from the cu13 toolkit (Built on Wed_Aug_20_01:53:56_PM_PDT_2025).

A few followups in case any of these matters for the difference between your working cu13 repro and our failing one:

1. Could you share the working setup from your Spark? Specifically Python version, manylinux tag for the libs-cu13 wheel you installed (cp310/cp311/cp312/cp313/cp314_..._aarch64.whl), and CUDA runtime / driver versions. We're on cp312-cp312-manylinux_2_28_aarch64 for libs-cu13 and cp312-abi3-manylinux_2_28_aarch64 for jit-cache.
2. Is there a particular toolchain-side fix expected to land in CUDA 13.0.x (or a 13.1 release) that closes the gap? Or a separate nvidia-cutlass-dsl-libs-cu13 4.5.x with the fix that we should pull?
3. Anything specific about how the kernel is invoked from cute-dsl that would help us validate the build? Our failing path is FlashInfer's flashinfer_b12x MoE backend (Integrate flashinfer b12x MoE and FP4 GEMM kernels for SM120/121 vllm-project/vllm#40082, my earlier comment there has the full launch command). The actual emission site is MmaSM120BlockScaledOp → nvvm.mma.block_scale {aType=e2m1, bType=e2m1, blockScaleFormat=ue4m3, scaleVecSize=x4, shape=<m=16,n=8,k=64>, cType=f32} with chip="sm_121a" (and the same failure also reproduces with chip="sm_120f", so it's not strictly sm_121a-specific from our end).

I can dump the full IR module + the PTX cute-dsl handed to ptxas if useful — CUTE_DSL_KEEP=ptx,ir-debug is set in the recipe so the artifacts should still be on the worker.

[depaulmillz]

In your docker container example above it looks like you are directly storing from the wheel. Any change you can pip install nvidia-cutlass-dsl[cu13] instead?

[idonati]

Update — your wheel-collision hint was the issue. The _mma ptxas error is gone with a properly-ordered install.

Two notes that may help future folks chasing the same bug:

Root cause (worth landing as a wheel-side fix?)

nvidia-cutlass-dsl-libs-base==4.5.0 and nvidia-cutlass-dsl-libs-cu13==4.5.0 both ship a file at the exact same path inside the package — 168 overlapping .py files including cute/nvgpu/warp/mma.py — but with different content / different SHA256:

$ unzip -p libs-base.whl '*RECORD' | grep 'warp/mma.py'
nvidia_cutlass_dsl/python_packages/cutlass/cute/nvgpu/warp/mma.py,sha256=T8iJjTYINVx5TKPJ0UsqcEpcYgQaQCvcOogVlPZgCoI,15232

$ unzip -p libs-cu13.whl '*RECORD' | grep 'warp/mma.py'
nvidia_cutlass_dsl/python_packages/cutlass/cute/nvgpu/warp/mma.py,sha256=yRcZZKWko3FxYxQop-OHHNVnl1Vyn7p0DwrkdeXG8yk,13904

The two versions are structurally different code paths:

• libs-base version: admissible_archs = ["sm_120a", "sm_121a"] (string literals), arch = CuTeDSL._get_dsl().get_arch_enum()
• libs-cu13 version: admissible_archs = [Arch.sm_120a, Arch.sm_121a] (enum), arch = BaseDSL._get_dsl().get_arch_enum()

When pip install nvidia-cutlass-dsl[cu13] is run from a clean slate, pip resolves order so cu13 wins. But when a previous install of either wheel is present (e.g., from pip install --force-reinstall of just one wheel, or from a Dockerfile layer that installed only base earlier), the file on disk can be the libs-base version even after a normal upgrade — and that variant trips the _mma ptxas error every time.

Concrete proof: the bytes-on-disk SHA256 of mma.py differs between a v6 image (force-reinstalled wheels, layered on top of an older base that had cutlass-dsl 4.4.2):

v6 / failing:   e49c79fd43e28b00442b10c294277702ad3fc0c973c72b84cd00e3dcfab14f48
v7 / working:   c9171964a5a4a37171631428a7e3871cd5679755729fba740f0ae475e5c6f329

After explicit pip uninstall -y nvidia-cutlass-dsl* && rm -rf /usr/local/lib/python3.12/dist-packages/nvidia_cutlass_dsl* && pip install --upgrade "nvidia-cutlass-dsl[cu13]>=4.5.0" (single RUN, exactly what you suggested), the file on disk is the cu13 SHA256, and the _mma ptxas error is gone.

Suggestion: make libs-base not ship the duplicate copy of any file that libs-cu* re-ships, or make libs-cu* list its overlapping files in the wheel RECORD with a unique marker so a later pip install libs-base doesn't silently revert them. This will save the next person several days. Probably worth a separate issue against the packaging.

New downstream error (just dropping here for the record, will pursue elsewhere)

With the right mma.py in place, the kernel compiles cleanly and we make it through to the first MoE forward pass during warmup. Hit a different error there — looks like an EP-slicing shape mismatch in the FlashInfer b12x kernel call rather than a cutlass-dsl issue:

ValueError: Mismatched b_w13.shape[2] on argument #9 when calling: `__call__(
  a_input: Tensor([8192, 6144], bfloat16),
  ...
  packed_a:    Tensor([65536, 3072, 48], float4_e2m1fnx2),      # <-- 48 = local experts/rank
  ...
  b_w13:       Tensor([4096, 3072, 384], float4_e2m1fnx2),      # <-- 384 = global experts
  b_down:      Tensor([6144, 1024, 384], float4_e2m1fnx2),      # <-- 384 global
  row_counts:  Tensor([48], int32),                             # <-- 48 local
  weight_expert_ids:      Tensor([48], int32),                  # <-- 48 local
  global_to_local_expert: Tensor([384], int32),
  input_global_scale:     Tensor([384], float32),
  alpha:                  Tensor([384], float32),
  ...
)`, expected to be 384

384 total experts / TP=EP=8 = 48 local experts/rank, which matches packed_a's last dim and row_counts. But b_w13 / b_down weights are still at the global 384 (not sliced to 48). I'll take this one to FlashInfer / vLLM as a separate issue — just thought you'd want closure on the original report. Thank you for the fast and precise debugging help!

I'll close this once the wheel-collision avenue gets a separate tracking issue (or just confirm here if you'd like to keep it here).