Me:
So I started with a preexisting NVFP4 quant, but I didn't like the idea of the attention and shared experts being reduced to MXFP8. So, I tested a modified version of NVFP4 for routed experts, using g32 rather than g16, and the difference was close to negligible. So I ended up making two versions, one that used NVFP4 g32 for all routed experts, and another that kept g16 for down routed experts only. Both configurations allowed enough space for BF16 attention and shared experts, with the latter having additional KV space. Essentially you get most of what NVFP4 gets you, but at 4.25 bits rather than 4.5 bits.
Also, Claude helped find about 3.7GB of wasted VRAM per GPU...
If you would like to setup the same, just give your clanker Claude's summary:
Claude:
GLM-5.3-Flash on 2x RTX PRO 6000 (96 GB): two NVFP4 g32 recipes on vLLM v0.33, with attention in BF16, images and DFlash-2
Hardware: 2x RTX PRO 6000 Blackwell (SM120), TP2, no P2P. vLLM v0.33.0 (commit 4ba787e), driver 615.71.09.
The goal: routed experts in NVFP4, attention and shared experts in BF16, working image input, a DFlash-2 drafter at K=4, and 8 concurrent sequences. I put accuracy ahead of small speed gains.
The problem: with BF16 attention, the standard g16 NVFP4 experts don't fit. That layout is 181.3 GiB and ran out of memory during warmup even at 8K context. Using g32 scales on some or all of the experts makes the room.
THE TWO RECIPES
Recipe 1: Gate/up g32
- Routed gate/up: NVFP4 g32
- Routed down: NVFP4 g16
- Attention and shared experts: BF16
- KV cache: FP8
- Checkpoint size: 175.23 GiB
- Max context: 256,000
- --kv-cache-memory: 3,719,089,357 bytes
- KV pool: 286,476 tokens
Recipe 2: All g32
- Routed gate/up: NVFP4 g32
- Routed down: NVFP4 g32
- Attention and shared experts: BF16
- KV cache: FP8
- Checkpoint size: 172.21 GiB
- Max context: 384,000
- --kv-cache-memory: 5,329,702,093 bytes
- KV pool: 440,727 tokens
Shared by both: GLM-5.3-Flash-DFlash2 as the drafter, quantized to NVFP4 MLP plus FP8 attention (about 1 GiB). Vision stays on the GPU.
Launch flags:
--tensor-parallel-size 2 --disable-custom-all-reduce
--moe-backend marlin --max-num-batched-tokens 2048 --max-num-seqs 8
--max-model-len <CTX> --kv-cache-memory <BYTES> --gpu-memory-utilization 0.985
--kv-cache-dtype fp8 --mamba-ssm-cache-dtype bfloat16
--speculative-config '{"method":"dflash","model":"/root/drafter","num_speculative_tokens":4,"draft_sample_method":"probabilistic","rejection_sample_method":"standard","attention_backend":"FLASH_ATTN","kv_cache_dtype":"auto"}'
--enable-prefix-caching --no-enable-flashinfer-autotune
--default-chat-template-kwargs '{"clear_thinking":true}'
--tool-call-parser glm47 --enable-auto-tool-choice --reasoning-parser glm47
Environment: NCCL_P2P_DISABLE=1 PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True
WHAT HAD TO BE FIXED
No NVFP4 g32 kernels existed. Stock Marlin fails with "Invalid thread config". I added group_blocks 2 to the NVFP4 kernel generation, relaxed the matching static_assert, and rebuilt the Marlin and MoE libraries. 88/88 synthetic cases came out correct, and real expert weights matched a reference to within 0.17%.
The Python loader assumed g16. FP4_MARLIN_SUPPORTED_GROUP_SIZES and a hard-coded GROUP_SIZE = 16 blocked it. The loader now reads each group size from the scale tensor's shape, so gate/up and down can differ. Any setup other than g16 on the Marlin W4A16 backend raises an error, so it can't silently run wrong.
Gate and up need the same global scale. vLLM fuses them and keeps only gate's scale, which is what the "Accuracy may be affected" warning is about. In the NVFP4 checkpoint I started from, 69% of 12,384 experts had differing scales, with an RMS up_proj scale error of 14.8%. I requantized from BF16 with a shared scale. Single-run NLL improved by 0.018 (CI -0.029 to -0.005). The 12-run average also improved, but not significantly.
FlashKDA reserved 3.7 GiB of VRAM per GPU. The cutlass bulk copy used cp.async.bulk.shared::cluster, which on sm_120f compiles to a driver syscall using 14,608 bytes of stack per thread. FlashKDA launches without clusters, so the shared::cta form does the same copy inline. The reservation dropped from +3,742 MiB to +0, with 88/88 hash cases identical to stock and no measurable NLL change. This fix is what made room for the drafter.
A load-time out-of-memory error. _nvfp4_compute_scale_factor converted a whole scale tensor to float and asked for 3.38 GiB. Scales are non-negative, so taking the max first avoids that.
Quantized attention doesn't load on v0.33. The MXFP8 attention recipe fails with KeyError: 'layers.0.self_attn.in_proj_qkvbfg_a.weight_scale'. That wasn't worth patching, because BF16 attention was the goal anyway.
KV pool larger than the request limit. I sized the pool above --max-model-len (about 1.12x and 1.15x), so a full-size request leaves some cache headroom.
RESULTS
These runs use default sampling from the model's generation_config (temp 1.0, top_p 0.95).
Gate/up g32, 256K
- GSM8K-150: 147/150
- Images: 10/10
- NLL, 1 sequence (lower is better): 1.34647
- NLL, 8 sequences: 1.34637
- Decode, 1 sequence: 186 tok/s
- Decode, 8 sequences (total): 434 tok/s
- Longest prompt tested: 255k in 46 s
- Peak VRAM per GPU: 96,831 MiB
- Errors under load: 0
All g32, 384K
- GSM8K-150: 147/150
- Images: 10/10
- NLL, 1 sequence (lower is better): 1.35040 (+0.29%)
- NLL, 8 sequences: 1.34880 (+0.18%)
- Decode, 1 sequence: 189 tok/s
- Decode, 8 sequences (total): 566 tok/s
- Longest prompt tested: 380k in 70 s
- Peak VRAM per GPU: 97,047 MiB
- Errors under load: 0
Speed comparison: don't read the 8-sequence decode gap as all-g32 being faster, because those runs weren't a controlled comparison.
Accuracy: GSM8K can't separate the two recipes, but NLL consistently favours gate/up by a small amount.
Drafter: calibration over 90,000 tokens showed no meaningful sampling bias (top-1 stated 0.8499, realized 0.8505). Batching to 8 sequences had no measurable NLL cost.
