Report: NVFP4 kernel selection on SM120/SM121 (results, no code changes)

[danielhanchen]

Results only, no code changes.

Update 2026-09-08: bug 1 has been fixed upstream. vllm#55170 merged 2026-09-08 as 13cf9e05c and moves the weight only kernel below all three W4A4 kernels. The workaround below is still required on every build that predates that commit, which includes every tag up to and including v0.29.0. Bug 2 is unchanged and still live on main.

On SM120/SM121 (RTX 50, RTX PRO 6000, DGX Spark), vLLM does not use the W4A4 kernels for our NVFP4 quants by default. Two upstream bugs send them down slower paths. Use these flags:

export VLLM_DISABLED_KERNELS=FlashInferCuteDslNvFp4W4A16LinearKernel
vllm serve <nvfp4-model> --moe-backend b12x    # needs: pip install b12x

Verified effect on kernel choice:

unset                                                          -> FlashInferCuteDslNvFp4W4A16LinearKernel   (dequant to BF16)
VLLM_DISABLED_KERNELS=FlashInferCuteDslNvFp4W4A16LinearKernel  -> FlashInferCutlassNvFp4LinearKernel        (native FP4)

The two bugs

	Bug	Cause	In a release?	Fixed?
1	Dense NVFP4 layers pick the weight only kernel over the native W4A4 kernels that do run on SM12x	882ca8d696 (vllm#53014, 2026-08-31) added the kernel at position 2 of _POSSIBLE_NVFP4_KERNELS[CUDA] with an is_supported() accepting sm_12x. First match wins. The same PR gates its explicit weight only branch to sm_100/103, so the auto list entry contradicts its own intent	No. Main and nightly only. v0.28.0, v0.29.0rc4 and v0.29.0 do not contain the class	Yes, #55170 merged 2026-09-08 (13cf9e05c), W4A16 kernel moved to position 5 below all three W4A4 kernels. Not in any tag yet
2	B12X MoE backend is never auto selected, falls back to Marlin	2740c817ff (vllm#52018) added the enum, kernel mapping and alias but never added B12X to AVAILABLE_BACKENDS	Yes, from v0.28.1rc0 onward, including v0.29.0. Not in v0.28.0, which predates the backend	No. Still absent from AVAILABLE_BACKENDS on main

Upstream state as of 2026-09-08:

Link	Bug	State
#55397	1	Open, but the defect it reports is fixed. Stale open, nobody closed it
#55170	1	Merged 2026-09-08. This is the fix
#55405	1	Open but superseded. Called a duplicate of #55170 by its author on 2026-09-07 and merge conflicted since
#54666	2	Open, no upstream traction
#54669	2	Open, no human review
#54719	2	Open, no human review. Competing duplicate of #54669

Note FLASHINFER_B12X is excluded from auto selection deliberately and is documented as such, though only in a source comment. Do not force it, it crashes with an illegal memory access on SM121. Plain B12X is the one that should be selected.

SM120 versus SM121, and a divergence worth recording

vLLM treats sm_120 and sm_121 as one family throughout. There is no exact sm_121 test anywhere in the NVFP4 selection path; the deciding gate is Platform.is_device_capability_family, which is literally

return (current_capability.to_int() // 10) == (capability // 10)

That is the direct analogue of blackwell_mma_available(cc) on the llama.cpp side, which also covered all of sm_12x. The two stacks have now diverged. In unslothai/llama.cpp#198 that gate was narrowed to an exact equality test on GGML_CUDA_CC_DGX_SPARK (1210), so it is sm_121 only. vLLM has not been narrowed and still catches every sm_12x part. A gate transplanted from one stack to the other is therefore no longer equivalent, and "SM12x" means different sets of hardware in the two codebases.

One place where the family wide treatment is genuinely wrong, below vLLM: of the native W4A4 kernels, FlashInferCutlassNvFp4LinearKernel and CutlassNvFp4LinearKernel run on both archs, but FlashInferB12xNvFp4LinearKernel is gated on has_device_capability(120) and has_flashinfer_b12x_gemm(), and has_flashinfer_b12x_gemm() is a bare hasattr probe with no arch check, while FlashInfer's own dense_blockscaled_gemm_sm120.py raises on anything that is not sm_120 and its release wheels build 120a/120f and never 121a. So that kernel is sm_120 only in practice despite a family wide gate on the vLLM side. It is harmless today only because FlashInferCutlassNvFp4LinearKernel outranks it on both archs. This is the same root cause the AVAILABLE_BACKENDS comment names, an SM121 CUTLASS guard being used to justify an exclusion applied to sm_120 as well.

Bug 1 itself is real on both archs. It was measured here on GB10 (sm_121); upstream #55170 independently measured the same regression on an RTX PRO 6000 Blackwell Max-Q (sm_120), same model, auto at TTFT 1455.2 ms / 74.81 tok/s against flashinfer_cutlass at 926.1 ms / 92.63 tok/s. Bug 2 is arch independent: AVAILABLE_BACKENDS is a flat list with no capability condition, so B12X is missing everywhere.

Why it matters

Per GEMM, 27B gate_up (N=34816 K=5120), CUDA graph replay. Time relative to W4A4, so higher is slower:

M	W4A16 CuTe DSL	W4A16 Marlin	W8A8 FP8	BF16
1	0.92	0.93	1.42	3.05
16	0.98	0.94	1.47	2.77
64	1.76	1.36	1.83	3.52
256	3.06	3.23	2.17	3.28
1024	4.15	4.63	1.73	2.99
4096	10.00	5.34	2.14	3.17

Peak observed: W4A4 80 to 148 TFLOPS, W8A8 45 to 55, BF16 30 to 33, Marlin 25 to 35, W4A16 CuTe DSL 20 to 22 and dropping to 8 to 10 at M=4096.

So the loss is in prefill and large batches, not single stream decode. Matches the upstream report of +27.5% prefill (pp2048 c1, 1423 to 1814 tok/s) on 2x DGX Spark with the kernel disabled.

SASS confirms it, all sm_121a: W4A4 runs OMMA.SF.16864.F32.E2M1.E2M1.UE4M3.4X (native block scaled FP4), W4A16 runs F2FP.F16.E2M1.UNPACK_B then HMMA.16816.F32 (unpack, then FP16 tensor cores), and 936 Marlin instantiations contain zero FP4 instructions. GB10 does expose FP4 MMA via mma.sync, the "no FP4 tensor cores" claim refers to tcgen05 on sm_100. That holds on sm_120a as well as sm_121a, the two differ only in which PTX ISA version first listed them (8.7 for sm_120a, 8.8 for sm_121a).

Caveat

My end to end serving runs on a battery throttled laptop were inconclusive, all configs within 10% and variance larger than the differences. The per GEMM table is the reliable measurement. The MoE model also masks the effect, since with stock flags the experts already get FLASHINFER_CUTLASS, a real W4A4 kernel, and only the small dense projections fall back. Bug 1 matters most for dense models.

Other hiccups worth recording

• Multimodal Qwen3.x checkpoints need --limit-mm-per-prompt '{"image":0,"video":0}' or startup profiling crashes.
• HF_HUB_DISABLE_XET=1 when pulling large quants on Spark, the Xet path thrashed the box, plain path holds about 90 MB/s.
• Under WSL2 the GPU allocatable memory is capped by the VM RAM in .wslconfig, not the host unified pool.
• FlashInfer JIT takes about 43 minutes on first start without a prebuilt aarch64 flashinfer_jit_cache with sm_121a cubins.

Tested on GB10 (sm_121), WSL2 Ubuntu 24.04, CUDA 13.0, vLLM 0.28.1rc1.dev345, torch 2.13.0+cu130, FlashInfer 0.6.18, with unsloth/Qwen3.8-27B-NVFP4 and unsloth/Qwen3.6-35B-A3B-NVFP4-Fast.

[chatgpt-codex-connector]

Codex Review Summary

This comment shows the latest Codex review activity on this pull request.

Review	Status	Commit	Review trigger
📝 Code Review	✅ Completed 2026-09-09T06:46:27.843589Z	7ed7d05	Manual request

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

• Open a pull request for review
• Mark a draft as ready
• Comment "@codex review" or "@codex security review".

Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings.

[danielhanchen]

@codex security review

[chatgpt-codex-connector]

Codex Review: Didn't find any major issues. Swish!

Reviewed commit: 7ed7d058eb

ℹ️ About Codex in GitHub

Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you

• Open a pull request for review
• Mark a draft as ready
• Comment "@codex review".

If Codex has suggestions, it will comment; otherwise it will react with 👍.

Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".