FlashInfer CUTLASS MoE fails TMA descriptor init on SM121

[TensorRT-LLM][ERROR] Assertion failed: Failed to initialize cutlass TMA WS grouped gemm. Error: Error Internal

Symptom

The day after fixing the FP8 arch-guard trap, the same model — mistralai/Mistral-Small-4-119B-2603-NVFP4 — ran clean for a long time and then died a different way. The server had been up since roughly 05:32 UTC and handled traffic without issue for about four and a half hours before crashing at 10:08 UTC, mid-MoE-forward pass, on a request with 36,308 computed tokens and a single output token.

It didn't crash without warning. The lead-up shows repeated TMA descriptor initialization failures:

Error: Failed to initialize the TMA descriptor 715

This is a different subsystem from the FP8 scaled-mm trap — this failure is in the fused-MoE grouped-GEMM path, on a server whose MoE backend had resolved to FLASHINFER_CUTLASS.

Root cause

FlashInfer's fused-MoE kernels are JIT-compiled the first time they're needed and cached on disk as fused_moe_120.so, under ~/.cache/flashinfer/jit_cache/. The 120 in that filename is the tell: the kernel templates behind it target cutlass::arch::Sm120. Dispatch happens through get_cutlass_fused_moe_module(device_arch), and on SM121 that function maps straight onto the SM120 module — there's no separate SM121 build to fall back to. That would be fine if the two architectures agreed on everything relevant. They don't: the TMA descriptor format differs between SM120 and SM121, and the SM120 module's descriptors fail to initialize against real SM121 hardware.

Two things narrow this down further, without fully explaining it. Sequential requests are fine — only concurrent batches trigger the failure, and neither investigation pins down why. And NVIDIA's own DGX Spark NVFP4 fix (PR #38423, commit b4a2f3ac3) does not resolve it. That patch changed quant_scales = None to quant_scales = [] and improved kernel detection on the vLLM side — a real fix for what it targets — but the crash survives it unchanged, because the fault was never in vLLM's dispatch logic. It's baked into a .so FlashInfer already built for the wrong chip. vLLM can change every line of its own source and this still crashes, since vLLM never gets a chance to generate the kernel in the first place.

Fix

The working option for this model was --moe-backend cutlass — vLLM's own CUTLASS kernels, compiled from source with 12.1a, instead of FlashInfer's precompiled 120 build. That held for 9+ hours under load. Two backends to avoid for this specific failure: marlin (FP4→FP16 decompression — very slow) and flashinfer_cutlass / auto (the same TMA crash).

That is specific to this model, not a general SM121 answer — NVIDIA's own DGX Spark guidance points at marlin, and a different NVFP4 MoE model needed a completely different setting to get stable. Before copying this flag onto anything other than Mistral Small 4, see the follow-up on which MoE backend actually works.