perf(linear): run nvfp4 linears on nvfp4 activations on sm_120

[jason-fxz]

NVFP4 linears on sm_120 for checkpoints whose linears carry static activation scales (ModelOpt NVFP4, compressed-tensors NVFP4 with dynamic: local inputs): prefill runs W4A4 GEMMs on NVFP4 activations, decode-sized batches run new W4A16 kernels on the same row-major weight. One weight copy, no repack; W4A16 checkpoints are unchanged.

Changes

• Two W4A4 NVFP4 linear kernels, first in the table on sm_120 with Gluon: b12x (FlashInfer fp4_quantize + mm_fp4(backend="b12x"), picked when flashinfer is installed) and cublas (our activation quant + torch scaled_mm, cuBLASLt). Both keep the checkpoint's row-major [N, K/2] codes with 128x4-swizzled block scales and take one weight global per tensor: a fused weight uses its largest per-part global with the block scales unchanged, as vLLM does, and a load-time warning names the case when the parts differ by more than 5%. --quant-backend linear.nvfp4=triton keeps main's W4A16 with each part's global.
• Merge after fix(muse_glimmer)!: keep the attention gate out of the qkv fusion #741, which keeps Muse-Glimmer's attention gate out of the q/k/v fusion: Muse-Glimmer-30B-NVFP4 quantizes that gate with a different global than q/k/v (down to 0.40x), which one tensor-wise global would scale wrongly. With it, every fused NVFP4 weight of the local NVFP4 checkpoints (Qwen3.6-27B, Qwen3.8-27B, Gemma-4-31B, Muse-Glimmer-30B, GLM-4.7) has equal per-part globals.
• kernel/gluon/nvfp4/linear/sm120_gemv.py: W4A16 kernels for small batches on that weight. gemv_fma (CUDA cores, a warp owns whole rows, cvt.rn.f16x2.e2m1x2 decode, T = 1-2) and gemv_mma (bf16 mma.sync, weight rows on M and tokens on N, each warp a K slice; a fixed permutation of k inside each 128-k block makes plain 16-byte code loads and 64-byte activation loads the mma fragments). K is a constexpr and the K loop is unrolled. gemv_config picks the kernel and tile. b12x routes T <= 8 to them, cublas T <= 16.
• The activation quant kernel splits rows into blocks (fixes the slow large-K case) and can write the 128x4-swizzled scales scaled_mm / mm_fp4 take.
• Tests: both linears against a torch W4A4 / W4A16 emulation across the routes, row-strided and empty inputs, every gemv_config tile against a dequant reference, and the swizzled scale layout. b12x refuses N % 8 != 0 (flashinfer returns wrong output there) and checks flashinfer's own CUDA toolkit version.

Cache policy: the activation loads use evict_last; the weight keeps the default policy. evict_first on the weight made these kernels faster alone but, in a model step, left earlier dirty L2 lines (GDN state) for the next layer's attention-side kernels to write back (nsys: FP8 GEMMs / GDN update / attention +12%, net slower; a CUDA microbench with ncu DRAM byte counts shows the same write-back moving from the stream to the kernels after it).

Results

Against main's W4A16 linear (--quant-backend linear.nvfp4=triton, the same code as main). torch 2.14.1+cu130, Triton 3.8.0, flashinfer 0.6.18.

End to end (LLM, CUDA graphs, naive cache; TTFT at ~1k / 4k / 8k prompt tokens; decode tok/s over 256 tokens at bs concurrent requests, median of 3):

RTX PRO 6000 (driver 595.91)

model	kernel	TTFT 1k	4k	8k	decode bs=1	bs=4	bs=16
Qwen3.8-27B-NVFP4	b12x	2.13x	1.95x	1.89x	-0.4%	+3.7%	+0.2%
	cublas	2.16x	1.94x	1.87x	+0.2%	+3.8%	-2.0%
Gemma-4-31B-IT-NVFP4	b12x	2.17x	1.79x	1.63x	+0.3%	+6.7%	+2.8% *
	cublas	2.13x	1.76x	1.62x	+0.3%	+6.6%	-0.6% *

* radix cache with unique prompts: Gemma-4-31B cannot allocate its KV pool at bs=16 under the naive cache.

RTX 5090 (driver 580.105.08), Qwen3.8-27B-NVFP4 (Gemma-4-31B does not fit):

kernel	TTFT 1k	4k	8k	decode bs=1	bs=4	bs=16
b12x	2.62x	2.44x	2.40x	+0.9%	+9.4%	+5.5%
cublas	2.69x	2.43x	2.38x	+0.8%	+9.3%	0.0%

Decode linears: each model's NVFP4 linears stacked to ~2 GB of weights (every layer reads from DRAM), one CUDA graph, through finalize + apply; speedup over main, b12x / cublas:

GPU	model	T=1	T=2	T=4	T=8	T=16	T=32
RTX PRO 6000	Qwen3.8-27B	1.02 / 1.02	1.11 / 1.11	1.10 / 1.10	1.09 / 1.10	1.01 / 0.95	1.09 / 0.79
	Gemma-4-31B	1.01 / 1.02	1.22 / 1.22	1.20 / 1.20	1.21 / 1.19	1.12 / 0.94	1.02 / 0.76
	Muse-Glimmer-30B	1.02 / 1.02	1.20 / 1.20	1.20 / 1.20	1.18 / 1.17	1.13 / 0.94	1.18 / 0.85
	GLM-4.7 (dense layer)	1.00 / 1.00	1.06 / 1.06	1.06 / 1.06	1.05 / 1.05	0.93 / 0.97	1.05 / 0.71
RTX 5090	Qwen3.8-27B	1.02 / 1.01	1.26 / 1.26	1.26 / 1.26	1.21 / 1.21	1.20 / 1.04	1.25 / 0.91
	Gemma-4-31B	1.01 / 1.01	1.20 / 1.20	1.18 / 1.18	1.14 / 1.14	1.13 / 0.98	1.11 / 0.84
	Muse-Glimmer-30B	1.00 / 1.00	1.19 / 1.19	1.19 / 1.18	1.15 / 1.14	1.14 / 0.93	1.36 / 0.99
	GLM-4.7 (dense layer)	1.02 / 1.02	1.09 / 1.09	1.08 / 1.08	1.05 / 1.06	1.00 / 0.94	1.19 / 0.83

The Muse-Glimmer-30B rows were measured with the gate fused into qkvg_proj.

From T = 48 main's W4A16 drops a tier and both W4A4 kernels run 1.6-2.8x faster. cublas trails main from T ~ 17 to ~ 40 (cuBLAS runs one 128x128 tile shape at every M): 0.71-0.91x at T = 32 above, -2.0% / -0.6% at bs=16 end to end.

In that measurement the small-batch kernels read the weight at 87-96% of the DRAM read bandwidth for T <= 8 (RTX PRO 6000, 1530 GB/s measured), at or above ninfer's NVFP4 linear in the same measurement (Qwen3.8-27B MLP, best of its A16 / A4 paths: 102.7 / 105.4 / 108.3 vs 104.5 / 108.7 / 116.3 us per layer at T = 1 / 4 / 8).

Test plan

• pytest tests/kernels/test_nvfp4_w4a4_linear.py on RTX PRO 6000 (and on RTX 5090 before the per-part global change): 56 passed, 20 skipped (N = 12, refused by both kernels).
• pytest tests/moe/test_nvfp4_backends.py tests/models/test_quant_config.py tests/kernels/test_e4m3_compat.py: 5 failures in test_quant_config.py, the same at the base commit (they come from the local HF cache).
• Greedy outputs, Muse-Glimmer-30B-NVFP4 on top of fix(muse_glimmer)!: keep the attention gate out of the qkv fusion #741 (RTX PRO 6000): b12x gives the same text as main's W4A16 (linear.nvfp4=triton) on a three-sentence prompt.
• Greedy outputs of the benchmark prompt: Qwen3.8-27B-NVFP4 identical to main in the first 100 characters for both kernels; Gemma-4-31B-IT-NVFP4 differs from main (W4A4 prefill changes the numerics) on a degenerate repeated-text prompt where main's output is repetitive too.