SM120 (Bug) (With FIx)(RTX Blackwell) NVFP4 MoE: CUTLASS Grouped GEMM Produces Garbage Output; Fixed via FlashInfer SM120 Patches + compute_120f (CUDA 13.0) — 39 tok/s Native FP4

[brandonmmusic-max]

NVFP4 MoE on SM120 (RTX PRO 6000 Blackwell): Full Debug Report

Title

CUTLASS & FlashInfer NVFP4 MoE Grouped GEMM Fails on SM120 Desktop Blackwell GPUs — Debug Journey, Patches, and Benchmark Results

All native FP4 MoE backends produce garbage output or crash on SM120 (compute_120) due to broken CUTLASS grouped GEMM templates. Through systematic patching of FlashInfer 0.6.5's SM120 capability checks and CuTe DSL architecture restrictions, we achieved the first known correct native FP4 MoE output on desktop Blackwell — albeit at reduced speed (14.6 tok/s vs Marlin's 46-49 tok/s) due to FlashInfer autotuner falling back to slow kernel tactics after TMA WS grouped GEMM initialization failures.

---

Environment

Component	Detail
GPUs	4x NVIDIA RTX PRO 6000 Blackwell Workstation Edition (96GB GDDR7 each, 384GB total)
Compute Capability	SM 12.0 (sm_120, NOT sm_120a)
Interconnect	PCIe (no NVLink)
Driver	582.16
OS	Windows 11 Pro + WSL2 Ubuntu 22.04
CUDA	12.8 (primary), 13.0 (available for JIT)
PyTorch	2.10.0+cu128
vLLM	0.17.0
FlashInfer	0.6.5 (upgraded from 0.6.4)
CUTLASS	4.2.1 (vendored in vLLM), 4.4.1 (tested separately)

Model

Parameter	Value
Model	nvidia/Qwen3.5-397B-A17B-NVFP4
Total Params	397B (17B active per token)
Experts	512 routed + 1 shared, 10 routed per token
Quantization	NVFP4 (FP4 weights with FP8 block scales)
Parallelism	TP=2 + PP=2 (optimal for PCIe)
KV Cache	FP8 e4m3
Max Seq Len	32,768

---

The Problem

NVFP4 MoE models produce garbage output (random whitespace, commas, fragments) on SM120 desktop Blackwell GPUs when using any backend that relies on CUTLASS grouped block-scaled FP4 GEMM kernels. Dense (non-MoE) FP4 GEMM works correctly — the issue is specifically in the grouped GEMM path used by MoE expert computations.

Symptom

Prompt: "What is the capital of Kentucky?"
Output: "  ,   ,  (!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!"

The model loads, serves requests, and generates tokens — but the MoE expert GEMM produces numerically wrong results, leading to incoherent output.

---

What We Tried (Chronological)

Phase 1: CUDA Kernel-Level Fixes (vLLM Source Rebuilds)

1. GDC (Grid Dependency Control) Barriers

• Hypothesis: Missing PDL synchronization barriers in CUTLASS grouped GEMM
• Action: Added -DCUTLASS_ENABLE_GDC_FOR_SM100=1 to CMakeLists.txt
• Finding: The flag was silently ignored! compute_120 (without a) doesn't define __CUDA_ARCH_FEAT_SM120_ALL, so the #ifndef CUTLASS_GDC_ENABLED guard evaluated to false
• Fix: Added -DCUTLASS_GDC_ENABLED directly as a compiler flag
• Result: GDC barriers now compiled as real PTX instructions (griddepcontrol.wait/launch), but still garbage output

2. FP32 Amax Computation

• Hypothesis: Half-precision amax in cvt_warp_fp16_to_fp4 causing quantization errors on SM120
• Action: Patched nvfp4_utils.cuh to compute per-block amax entirely in FP32 (fabsf/fmaxf instead of __habs2/__hmax2)
• Result: Still garbage. Scale computation was already FP32; the half-precision amax wasn't the root cause.

3. Pingpong Kernel Schedule

• Hypothesis: Cooperative schedule buggy on SM120, Pingpong might work
• Action: Changed SM120 GEMM from KernelScheduleAuto to KernelPtrArrayTmaWarpSpecializedPingpong
• Result: SEGFAULT. Pingpong schedule crashes on SM120.

4. compute_120a Architecture Flag

• Hypothesis: Desktop SM120 supports accelerated MMA instructions
• Action: Forced compute_120a gencode for FP4 kernel compilation
• Result: SEGFAULT. RTX PRO 6000 reports compute capability 12.0, not 12.0a. The a-specific instructions are not available on desktop Blackwell (confirmed by CUTLASS Issue SM 120 (Blackwell GeForce RTX 50 Series) Block-Scaled MMA Runtime Assertion Failure #2820).

5. CUTLASS 4.4.1 Upgrade

• Hypothesis: CUTLASS 4.4.1 changelog mentions SM120 fixes
• Action: Cloned CUTLASS 4.4.1, set VLLM_CUTLASS_SRC_DIR, rebuilt _C.abi3.so
• Critical Bug: First clone attempt silently got 4.2.1 due to CMake's FetchContent_Declare overwriting our clone with hardcoded GIT_TAG v4.2.1. Fixed by using VLLM_CUTLASS_SRC_DIR env var.
• Result: Still garbage. CUTLASS 4.4.1 has the same broken SM120 grouped block-scaled GEMM templates.

Phase 2: Alternative MoE Backends (FlashInfer)

vLLM supports 5 MoE backends for NVFP4:

1. VLLM_CUTLASS (default) — broken on SM120
2. FLASHINFER_TRTLLM — blocked by SM100-only capability checks
3. FLASHINFER_CUTLASS — blocked by SM120 capability checks + missing sm_120a in CuTe DSL
4. FLASHINFER_CUTEDSL — blocked by SM100-only capability checks
5. MARLIN — working W4A16 workaround (46-49 tok/s)

6. FlashInfer CUTLASS Backend (The Breakthrough)

Required patches (10+ files):

vLLM Capability Checks (3 files)

# trtllm_nvfp4_moe.py, flashinfer_trtllm_moe.py, flashinfer_cutedsl_moe.py
# Changed:
return p.is_cuda() and p.is_device_capability_family(100)
# To:
return p.is_cuda() and (p.is_device_capability_family(100) or p.is_device_capability_family(120))

FlashInfer JIT Architecture Filters (flashinfer/jit/fused_moe.py)

# Lines 62, 79, 238: Added major version 12
supported_major_versions=[10]      # -> [10, 12]
supported_major_versions=[10, 11]  # -> [10, 11, 12]

FlashInfer Compilation Context (flashinfer/compilation_context.py)

# Changed: major >= 9 adds "a" suffix (generates compute_120a which is needed for CUTLASS MMA)
# SM120 needs "a" suffix for MMA instructions, but not "f" (CUDA 13.0+ only)

CuTe DSL admissible_archs (5 files, 18+ locations)

flashinfer/data/cutlass/python/CuTeDSL/cutlass/cute/nvgpu/cpasync/copy.py  (4 locations)
flashinfer/data/cutlass/python/CuTeDSL/cutlass/cute/nvgpu/tcgen05/mma.py   (2 locations)
flashinfer/data/cutlass/python/CuTeDSL/cutlass/cute/nvgpu/tcgen05/copy.py  (3 locations)
flashinfer/data/cutlass/python/CuTeDSL/cutlass/cute/arch/mbar.py           (8 locations)
flashinfer/data/cutlass/python/CuTeDSL/cutlass/cute/arch/elect.py          (1 location)

Added "sm_120a" after every "sm_100a" in admissible_archs lists.

cuda.py Device Mapping

# Added:
(12, 0): ("Blackwell", "sm_120a", ["sm_120a"]),  # RTX PRO 6000

TRT-LLM C++ Launcher (flashinfer/data/csrc/trtllm_fused_moe_kernel_launcher.cu)

// Lines 417, 1345: Changed == to >=
TVM_FFI_ICHECK_EQ(major, 10)           // -> TVM_FFI_ICHECK_GE(major, 10)
TVM_FFI_ICHECK_EQ(std::get<0>(...), 10) // -> TVM_FFI_ICHECK_GE(...)

Additional Requirements

• nvcc must be in PATH (FlashInfer JIT needs it)
• FlashInfer JIT cache must be cleared after patching
• VLLM_NVFP4_GEMM_BACKEND=cutlass env var for dense layers (use vLLM native CUTLASS)

Result: CORRECT OUTPUT! First known native FP4 MoE on SM120 desktop Blackwell.

---

Benchmark Results

Launch Command (FlashInfer CUTLASS — Working Native FP4)

export PATH="/usr/local/cuda-12.8/bin:$PATH"  # or cuda-13.0 for compute_120f
export VLLM_NVFP4_GEMM_BACKEND=cutlass
export NCCL_CUMEM_ENABLE=0
export VLLM_WORKER_MULTIPROC_METHOD=spawn

python -m vllm.entrypoints.openai.api_server \
  --model nvidia/Qwen3.5-397B-A17B-NVFP4 \
  --dtype bfloat16 \
  --tensor-parallel-size 2 \
  --pipeline-parallel-size 2 \
  --max-model-len 32768 \
  --gpu-memory-utilization 0.92 \
  --trust-remote-code \
  --moe-backend flashinfer_cutlass

Speed Comparison

Backend	MoE Kernel	CUDA	Single User (tok/s)	4-User (per user)	Output
Marlin (--moe-backend marlin)	W4A16 dequant	12.8	46-49	~37	Correct
FlashInfer CUTLASS 120f	SM120 CUTLASS JIT	13.0	39.0	18.2	Correct
FlashInfer CUTLASS 120a	SM120 CUTLASS JIT	12.8	14.6-14.9	6.9-8.5	Correct
FlashInfer CUTLASS Hybrid	SM120 JIT + vLLM dense	12.8	14.8-14.9	6.9	Correct
vLLM Native CUTLASS	Grouped block-scaled	12.8	N/A	N/A	Garbage
CUTLASS 4.4.1 rebuild	Grouped block-scaled	12.8	N/A	N/A	Garbage
FlashInfer TRT-LLM	TRT-LLM cubins	12.8	N/A	N/A	Crash

Why FlashInfer CUTLASS is 3x Slower Than Marlin

FlashInfer's autotuner logs reveal the root cause:

flashinfer.jit: [Autotuner]: Skipping tactic <MoERunner> 14, due to failure:
[TensorRT-LLM][ERROR] Failed to initialize cutlass TMA WS grouped gemm.
Error: Error Internal (cutlass_kernel_file_gemm_grouped_sm120_M128_BS_group2.generated.cu:60)

All TMA warp-specialized grouped GEMM tactics fail to initialize on SM120 with compute_120a. The autotuner falls back to slower, non-TMA tactics. This is a CUTLASS template-level issue where SM120's TMA grouped GEMM doesn't work with the a suffix — it likely requires the f suffix (compute_120f) which is only available with CUDA 13.0+.

---

Key Technical Findings

1. compute_120 vs compute_120a vs compute_120f

Flag	CUDA Version	MMA Instructions	CUTLASS Grouped GEMM	Result
compute_120	12.8+	Not enabled	"Arch conditional MMA" error	Fails
compute_120a	12.8+	Enabled	TMA WS tactics fail, slow fallback	14.6 tok/s
compute_120f	13.0+ only	Full feature set	Potentially fast tactics	Testing

2. SM120 Desktop is NOT SM100 Compatible

Despite sharing the "Blackwell" brand, SM120 (desktop) and SM100 (datacenter) have different:

• Compute capability families (12 vs 10)
• Supported architecture features (a vs f suffix)
• Pre-compiled cubin compatibility (SM100 cubins crash on SM120)

3. The Broken Chain

vLLM CUTLASS grouped GEMM → garbage output (kernel correctness bug)
    ↓ upgrade CUTLASS 4.4.1
Still garbage (same templates, 0 SM120 changes)
    ↓ try FlashInfer CUTLASS
Blocked: SM120 not in capability checks
    ↓ patch 10+ files
Works with correct output, but slow (autotuner fallback)
    ↓ try FlashInfer TRT-LLM
Crash: hardcoded SM==10 in C++ + SM100-only cubins
    ↓ next: compute_120f with CUDA 13.0
Pending...

---

BREAKTHROUGH: compute_120f with CUDA 13.0

A DGX Spark (SM121) user achieved 35 tok/s with FlashInfer CUTLASS using 12.1f (CUDA 13.0). The f suffix enables the "full" SM120 feature set with working TMA WS grouped GEMM tactics.

Results: compute_120f Nearly Triples Speed

Metric	compute_120a (CUDA 12.8)	compute_120f (CUDA 13.0)	Marlin W4A16
Single user	14.6 tok/s	39.0 tok/s	46-49 tok/s
4-user concurrent	6.9 tok/s/user	18.2 tok/s/user	~37 tok/s/user

compute_120f enabled the fast TMA WS grouped GEMM tactics that failed with compute_120a. This confirms the f suffix is the correct architecture designation for SM120 desktop Blackwell GPUs.

Launch Command (CUDA 13.0 + compute_120f)

export PATH="/usr/local/cuda-13.0/bin:$PATH"
export VLLM_NVFP4_GEMM_BACKEND=cutlass
export NCCL_CUMEM_ENABLE=0
export VLLM_WORKER_MULTIPROC_METHOD=spawn

python -m vllm.entrypoints.openai.api_server \
  --model nvidia/Qwen3.5-397B-A17B-NVFP4 \
  --dtype bfloat16 \
  --tensor-parallel-size 2 \
  --pipeline-parallel-size 2 \
  --max-model-len 32768 \
  --gpu-memory-utilization 0.92 \
  --trust-remote-code \
  --moe-backend flashinfer_cutlass

Why 39 vs 49 tok/s?

The remaining ~20% gap vs Marlin is likely due to:

• FlashInfer CUTLASS autotuner may not select the absolute optimal tactic
• Native FP4 GEMM has activation quantization overhead (BF16 -> FP4 per-token)
• Further kernel tuning by FlashInfer team could close the gap
• Pipeline parallel bubble overhead affects native FP4 slightly differently than Marlin

---

Production Recommendation (Current)

Use Marlin for production until compute_120f results are confirmed:

python -m vllm.entrypoints.openai.api_server \
  --model nvidia/Qwen3.5-397B-A17B-NVFP4 \
  --dtype bfloat16 \
  --tensor-parallel-size 2 \
  --pipeline-parallel-size 2 \
  --moe-backend marlin \
  --max-model-len 32768 \
  --gpu-memory-utilization 0.95 \
  --trust-remote-code

Required env vars:

export NCCL_CUMEM_ENABLE=0
export VLLM_WORKER_MULTIPROC_METHOD=spawn

---

Related Issues

• CUTLASS #2820 — SM120 Block-Scaled MMA Runtime Assertion Failure
• CUTLASS #2800 — BlockScaledMmaOp restricts FP4 to sm_100a only
• vLLM #33416 — NVFP4 MoE kernels fail on RTX Blackwell (SM12.0)
• vLLM #33333 — FLASHINFER_CUTLASS not supported on SM120
• vLLM #31085 — Add SM120 support for native NVFP4 MoE kernels
• FlashInfer #2577 — mm_fp4 GEMM broken on SM120
• NVIDIA Forum — 35 TPS with FlashInfer 12.1f on DGX Spark

---

Files Patched (Complete List)

FlashInfer 0.6.5

File	Change
flashinfer/compilation_context.py	Arch suffix logic for SM120
flashinfer/jit/fused_moe.py (3 locations)	Added supported_major_versions 12
flashinfer/data/csrc/trtllm_fused_moe_kernel_launcher.cu (2 locations)	ICHECK_EQ -> ICHECK_GE
flashinfer/data/cutlass/python/CuTeDSL/cutlass/cute/nvgpu/cpasync/copy.py (4 locations)	Added sm_120a to admissible_archs
flashinfer/data/cutlass/python/CuTeDSL/cutlass/cute/nvgpu/tcgen05/mma.py (2 locations)	Added sm_120a to admissible_archs
flashinfer/data/cutlass/python/CuTeDSL/cutlass/cute/nvgpu/tcgen05/copy.py (3 locations)	Added sm_120a to admissible_archs
flashinfer/data/cutlass/python/CuTeDSL/cutlass/cute/arch/mbar.py (8 locations)	Added sm_120a to admissible_archs
flashinfer/data/cutlass/python/CuTeDSL/cutlass/cute/arch/elect.py (1 location)	Added sm_120a to admissible_archs
flashinfer/data/cutlass/python/CuTeDSL/base_dsl/runtime/cuda.py	Added (12, 0) device mapping

vLLM 0.17.0

File	Change
vllm/model_executor/layers/fused_moe/experts/trtllm_nvfp4_moe.py	Added is_device_capability_family(120)
vllm/model_executor/layers/fused_moe/flashinfer_trtllm_moe.py	Added is_device_capability_family(120)
vllm/model_executor/layers/fused_moe/flashinfer_cutedsl_moe.py	Added is_device_capability_family(120)

vLLM Source (CUDA kernel rebuilds — tested but not needed for FlashInfer path)

File	Change
vllm-src/CMakeLists.txt	Added -DCUTLASS_GDC_ENABLED, -DCUTLASS_ENABLE_GDC_FOR_SM100=1
vllm-src/csrc/quantization/fp4/nvfp4_utils.cuh	FP32 amax computation

---

Report date: March 8, 2026
Hardware: 4x RTX PRO 6000 Blackwell (SM120, 96GB each)
Tested by: Kentucky Local Counsel local inference team lead Brandon Music

[CristyNel]

I think my failure may be part of the same SM120 / NVFP4 / grouped-GEMM / fused-MoE family of issues.

Environment

• GPU: RTX 5090 / SM120
• CUDA: 13.0
• PyTorch: 2.10.0+cu130
• FlashInfer: nightly-v0.6.5-20260309
• vLLM: 0.17.0+cu130
• nvidia-cutlass-dsl: 4.4.1

Use case

I am trying to build FlashInfer AOT for:
NVIDIA-Nemotron-3-Nano-30B-A3B-NVFP4

with FlashInfer fused MoE / CUTLASS path enabled.

Failure point

AOT build fails while compiling:

moe_gemm_kernels_fp8_uint4.cu

from the fused_moe_120 path.

The failing compile line is in the family:

.../cached_ops/fused_moe_120/moe_gemm_kernels_fp8_uint4.cuda.o

with SM120 flags like:

-gencode=arch=compute_120,code=sm_120
-DCOMPILE_BLACKWELL_TMA_GEMMS
-DCOMPILE_BLACKWELL_SM120_TMA_GROUPED_GEMMS
-DUSING_OSS_CUTLASS_MOE_GEMM

Then Ninja stops with non-zero exit status during FlashInfer AOT.

Why I think it is related to #3096

Issue #3096 is already describing SM120 desktop Blackwell NVFP4 MoE grouped GEMM problems / broken output / FlashInfer SM120 patching requirements. My failure is earlier in the pipeline (AOT compile-time instead of runtime garbage output), but it is still in the SM120 fused-MoE grouped-GEMM path.

There is also a currently open FlashInfer issue (#2693) stating that Blackwell CuteDSL fused-MoE is broken on nvidia-cutlass-dsl >= 4.4.0 and works on 4.3.5. However, in my case vllm 0.17.0+cu130 requires nvidia-cutlass-dsl >= 4.4.0.dev1, so downgrading to 4.3.5 is not a clean option in the main runtime environment.

Question

Does this look like the same underlying SM120 grouped-GEMM / NVFP4 / CUTLASS-CuTeDSL issue family as #3096, just surfacing as a FlashInfer AOT compile failure at moe_gemm_kernels_fp8_uint4.cu instead of runtime corruption?

[brandonmmusic-max]

To me it looks like it's related (possibly the same thing). I won't pretend to be the expert of experts, but I'm thinking the kernal issue is upstream, so it's got the opportunity to trigger in multiple different way downstream. Brandon M. Music W. Jeffrey Scott, P.S.C. Pronouns He, Him, His P.O. Box 608 311 W. Main Street Grayson, KY 41143 Tel: 606-474-5194 Fax: 606-474-5196 Bar #: 93785 e-mail: ***@***.*** Pay your invoice or account bill here: https://www.paypal.com/paypalme/wjeffreyscottpsc?country.x=US&locale.x=en_US This message and all attachments are confidential and may be protected by the attorney-client attorney work-product or other privileges. Any unauthorized review, use, disclosure or distribution is prohibited. If you believe this message has been sent to you in error, please notify the sender by replying to this message and delete the message without disclosing it to anyone. Thank you.

…

On Tue, Mar 10, 2026 at 1:24 PM Cristinel Nelu Buliga < ***@***.***> wrote: *CristyNel* left a comment (NVIDIA/cutlass#3096) <#3096 (comment)> I think my failure may be part of the same SM120 / NVFP4 / grouped-GEMM / fused-MoE family of issues. Environment - GPU: RTX 5090 / SM120 - CUDA: 13.0 - PyTorch: 2.10.0+cu130 - FlashInfer: nightly-v0.6.5-20260309 - vLLM: 0.17.0+cu130 - nvidia-cutlass-dsl: 4.4.1 Use case I am trying to build FlashInfer AOT for: NVIDIA-Nemotron-3-Nano-30B-A3B-NVFP4 with FlashInfer fused MoE / CUTLASS path enabled. Failure point AOT build fails while compiling: moe_gemm_kernels_fp8_uint4.cu from the fused_moe_120 path. The failing compile line is in the family: .../cached_ops/fused_moe_120/moe_gemm_kernels_fp8_uint4.cuda.o with SM120 flags like: -gencode=arch=compute_120,code=sm_120 -DCOMPILE_BLACKWELL_TMA_GEMMS -DCOMPILE_BLACKWELL_SM120_TMA_GROUPED_GEMMS -DUSING_OSS_CUTLASS_MOE_GEMM Then Ninja stops with non-zero exit status during FlashInfer AOT. Why I think it is related to #3096 <#3096> Issue #3096 <#3096> is already describing SM120 desktop Blackwell NVFP4 MoE grouped GEMM problems / broken output / FlashInfer SM120 patching requirements. My failure is earlier in the pipeline (AOT compile-time instead of runtime garbage output), but it is still in the SM120 fused-MoE grouped-GEMM path. There is also a currently open FlashInfer issue (#2693 <#2693>) stating that Blackwell CuteDSL fused-MoE is broken on nvidia-cutlass-dsl >= 4.4.0 and works on 4.3.5. However, in my case vllm 0.17.0+cu130 requires nvidia-cutlass-dsl >= 4.4.0.dev1, so downgrading to 4.3.5 is not a clean option in the main runtime environment. Question Does this look like the same underlying SM120 grouped-GEMM / NVFP4 / CUTLASS-CuTeDSL issue family as #3096 <#3096>, just surfacing as a FlashInfer AOT compile failure at moe_gemm_kernels_fp8_uint4.cu instead of runtime corruption? — Reply to this email directly, view it on GitHub <#3096 (comment)>, or unsubscribe <https://github.com/notifications/unsubscribe-auth/B7RPGFHIQDBL53527YTQ7ND4QBFULAVCNFSM6AAAAACWLL7PRCVHI2DSMVQWIX3LMV43OSLTON2WKQ3PNVWWK3TUHM2DAMZTGE3DINJWGE> . You are receiving this because you authored the thread.Message ID: ***@***.***>

[brandonmmusic-max]

Complete SM120 NVFP4 MoE Benchmark Report (4x RTX PRO 6000)

Following up on this issue with a comprehensive benchmark of every available MoE backend, MTP configuration, Docker image, kernel patch, and CUDA version tested on 4x RTX PRO 6000 Blackwell Workstation Edition (SM 12.0, 96GB GDDR7 each, PCIe).

Model: nvidia/Qwen3.5-397B-A17B-NVFP4 (397B total, 17B active/token, 512 routed + 1 shared expert)

All Configurations Tested (Single User Decode, 500-1000 tokens)

#	Configuration	MoE Backend	TP	MTP	CUDA	tok/s	Notes
1	Marlin TP=2+PP=2	Marlin W4A16	2+PP2	No	12.8	49	Correct output
2	orthozany Docker + MTP=2	FlashInfer CUTLASS (120a)	4	Yes	13.0	40	80 TMA WS tactics skipped
3	orthozany Docker + MTP=2 (warmed)	FlashInfer CUTLASS (120a)	4	Yes	13.0	41	80 TMA WS tactics skipped
4	Festr Docker (120a, no patch)	FlashInfer CUTLASS (120a)	4	Yes	13.1	26	80 TMA WS tactics skipped
5	Festr Docker + 120f env var	FlashInfer CUTLASS (120f)	4	Yes	13.1	41	80 TMA WS tactics skipped
6	Festr Docker + GDC flags + 120f	FlashInfer CUTLASS (120f+GDC)	4	Yes	13.1	41	80 TMA WS tactics skipped
7	VLLM_CUTLASS (split backend)	vLLM native CUTLASS	4	Yes	13.1	~5	Garbage output
8	Marlin via --moe-backend	Marlin W4A16	4	Yes	13.1	CRASH	Unquantized MoE error
9	Marlin via FORCE_FP8_MARLIN	Marlin W4A16	4	Yes	13.1	44	Low MTP acceptance
10	Pip nightly + Marlin + MTP=2	Marlin W4A16	4	Yes	13.0	39-40	MTP hurts (-22%)
11	Pip nightly + Marlin + NO MTP	Marlin W4A16	4	No	13.0	50.5	Best config
12	SGLang 0.5.8	FlashInfer	4	—	—	NaN	NaN outputs on SM120
13	TP=4 default (no Marlin force)	FlashInfer CUTLASS	4	No	12.8	6-7	Garbage output
14	TP=2 + EP=2 Marlin	Marlin W4A16	2+EP2	No	—	1.4-2.6	Expert parallel catastrophic on PCIe

Key Finding: MTP Hurts Marlin Performance

MTP acceptance rates on Marlin were only 61-85% (vs 89% community baseline). Marlin's W4A16 dequantization produces numerically different results than native FP4, causing MTP draft heads (trained on native FP4 activations) to mispredict. Without MTP: 50.5 tok/s. With MTP=2: 39.6 tok/s (-22%).

Things That Did NOT Fix the TMA WS Initialization Failure

• compute_120a vs compute_120f (same 80 tactics skipped either way)
• GDC flags (-DCUTLASS_ENABLE_GDC_FOR_SM100=1 -DCUTLASS_ENABLE_GDC_FOR_SM90=1)
• CUTLASS 4.4.1 upgrade
• CUDA 13.0 → 13.1 → 13.2 (driver 595.71)
• FlashInfer PR [BUG] [CUTLASS] Integer-Gemm gives incorrect results #2716 (GDC barrier fix)
• FP32 scale precision patches
• BF16 → FP16 activation casting

Patches Applied (12 total)

FlashInfer: trtllm_fused_moe_kernel_launcher.cu SM check fix, compilation_context.py 120f suffix, jit/fused_moe.py major version 12 support (3 locations), GDC compile flags, CuTe DSL admissible_archs (5 files, 18 locations), device capability mapping.

vLLM: is_device_capability_family(120) added to 5 MoE backend selection files.

PRs submitted: FlashInfer #2725, vLLM #36453.

What Would Unlock Native NVFP4 Performance

TensorRT-LLM PR #5823 ("fix moe regression for sm120") was merged but has NOT propagated to FlashInfer. Once the TMA WS grouped GEMM initializes correctly on SM120, we expect 70-86 tok/s baseline (native NVFP4) and ~130 tok/s with MTP=2.

The SM121 (DGX Spark) result of 356 TFLOPS for NVFP4 MoE grouped GEMM proves the SM12x architecture CAN handle this workload — the issue is SM120-specific tile config/SMEM initialization in the CUTLASS templates.

Environment: vLLM 0.17.0rc1.dev202, FlashInfer 0.6.4, PyTorch 2.9.1+cu128, Driver 595.71, WSL2 Ubuntu 22.04

[brandonmmusic-max]

Root Cause Found: SMEM Overflow in SM120 Grouped GEMM

After tracing through the entire CUTLASS template chain, I've identified the definitive root cause of the SM120 NVFP4 MoE grouped GEMM initialization failure.

The Bug

StageCountAutoCarveout (computed via sm100_compute_stage_count_or_override_blockscaled in sm100_blockscaled_umma_builder.inl) does not account for alignas(1024) padding on smem_A and smem_B in SharedStorage::TensorStorage (sm120_blockscaled_mma_array_tma.hpp). This causes it to compute 3 pipeline stages when only 2 actually fit in SM120's 101,376-byte SMEM budget (vs SM100's 232,448 bytes).

Exact Failure Location

GemmUniversalAdapter::initialize()          [gemm_universal_adapter.h:330]
  -> cudaFuncSetAttribute(MaxDynamicSharedMemorySize, smem_size)
  -> FAILS: sizeof(SharedStorage) > 101,376 bytes (sm120_smem_capacity_bytes)
  -> returns Status::kErrorInternal

The initialize_workspace() call succeeds -- the failure is AFTER workspace init, in the cudaFuncSetAttribute call. The SM90 array kernel shell (sm90_gemm_array_tma_warpspecialized_cooperative.hpp) has no static_assert checking SMEM capacity against the target architecture, unlike the SM120 dense kernel which has static_assert(SharedStorageSize <= sm120_smem_capacity_bytes).

Verification: StageCount Sweep

I patched moe_gemm_tma_ws_launcher.inl:497 to override StageCountAutoCarveout with fixed StageCount<N> for SM120:

Stages	SMEM Fits?	Kernel Init	Decode Speed	Notes
2	Yes	Success	4.8 tok/s	Kernel runs but too few stages to hide latency
3	No	CRASH	N/A	Original bug (auto-computed, alignment overflow)
4	No	CRASH	N/A	SMEM overflow
5	No	CRASH	N/A	SMEM overflow
Marlin W4A16	N/A	N/A	50.5 tok/s	Current workaround (no native FP4)

Why StageCount<2> Is Slow

With only 2 pipeline stages, the TMA WS grouped GEMM cannot overlap memory loads with compute. The SM120 dense (non-grouped) FP4 kernel also uses StageCount<2> (fp4_gemm_template_sm120.h:259) and gets ~39 tok/s for dense GEMM. The grouped GEMM at 4.8 tok/s suggests additional overhead from the PtrArray dispatch and expert routing.

Second Bug: Tile Validation Mismatch

are_tile_shapes_supported_sm120() only allows K=128/256 tiles (128x128x128, 128x128x256, 128x256x128, 256x128x128). But the dispatch (SHAPE_CASE) and heuristic (get_candidate_configs_sm120) both try K=64 tiles (128x128x64, 128x256x64, 256x128x64). These silently fail the validation gate and never run.

K=64 tiles would fit 7-11 pipeline stages in 101 KiB, providing much better pipeline overlap:

Tile Shape	~KiB/Stage	Stages in 101 KiB
128x128x128	~17.5	5 (but alignas padding pushes to overflow at 3+)
128x128x256	~34	2 (this is what actually runs)
128x128x64	~8.6	11
128x256x64	~13	7
256x128x64	~13	7

What NVIDIA Needs to Fix

1. Add K=64 tiles to are_tile_shapes_supported_sm120() (they're already in dispatch + heuristic)
2. Fix StageCountAutoCarveout to account for alignas(1024) padding on shared memory buffers
3. Add static_assert(SharedStorageSize <= sm120_smem_capacity_bytes) in the SM90 array kernel shell
4. Consider SM120-specific smaller tiles (64x128) for even better pipelining

Template Infrastructure Is Complete

The CUTLASS template chain for SM120 grouped GEMM NVFP4 is fully correct:

• sm120_blockscaled_mma_builder.inl -> supports PtrArray (grouped GEMM)
• KernelPtrArrayTmaWarpSpecializedCooperativeBlockScaledSm120 -> correct kernel schedule
• Sm120PtrArrayTmaWarpSpecialized -> correct epilogue dispatch
• SM90 array kernel shell -> handles SM120 at __CUDA_ARCH_FEAT_SM120_ALL

The ONLY issue is SMEM budget + tile validation. The infrastructure works -- it just needs tile configs that fit in 101 KiB with enough stages for performance.

Hardware: 4x RTX PRO 6000 (SM 12.0, 96GB), CUDA 13.0, CUTLASS 4.4.1, FlashInfer 0.6.4, vLLM 0.16.1rc1

Currently testing: K=64 tile validation fix (adding 128x128x64, 128x256x64, 256x128x64 to are_tile_shapes_supported_sm120()). Will update with results.

[depaulmillz]

Did you try this after this commit 73c59c0? It should correct the stage counts.

[brandonmmusic-max]

looking at the timing of the commit, I believe this was before that. Brandon M. Music W. Jeffrey Scott, P.S.C. Pronouns He, Him, His P.O. Box 608 311 W. Main Street Grayson, KY 41143 Tel: 606-474-5194 Fax: 606-474-5196 Bar #: 93785 e-mail: ***@***.*** Pay your invoice or account bill here: https://www.paypal.com/paypalme/wjeffreyscottpsc?country.x=US&locale.x=en_US This message and all attachments are confidential and may be protected by the attorney-client attorney work-product or other privileges. Any unauthorized review, use, disclosure or distribution is prohibited. If you believe this message has been sent to you in error, please notify the sender by replying to this message and delete the message without disclosing it to anyone. Thank you.

…

On Fri, Mar 13, 2026 at 12:26 PM dePaul Miller ***@***.***> wrote: *depaulmillz* left a comment (NVIDIA/cutlass#3096) <#3096 (comment)> Did you try this after this commit 73c59c0 <73c59c0>? It should correct the stage counts. — Reply to this email directly, view it on GitHub <#3096 (comment)>, or unsubscribe <https://github.com/notifications/unsubscribe-auth/B7RPGFCUATQ4VKU3NXF6P534QQZCLAVCNFSM6AAAAACWLL7PRCVHI2DSMVQWIX3LMV43OSLTON2WKQ3PNVWWK3TUHM2DANJWGM3DINRZGA> . You are receiving this because you authored the thread.Message ID: ***@***.***>