830K downloads later: what quantizing Qwen3.8-27B to NVFP4 for one RTX 5090 actually taught me

Our NVFP4 build of Qwen3.8-27B crossed 830,987 downloads this month. Around the same time, three other people published NVFP4 versions of the same model — including NVIDIA, who invented the format.

So I put all of them on one RTX 5090 and measured.

The result surprised me, and it changed what I think the interesting problem is.

The short version

Same box, same engine, same session, same sampling. Quality is a statistical tie — all three confidence intervals overlap, and long-context reasoning is a dead heat. We’re 25% faster with 43% more context than NVIDIA’s official build.

None of that came from a better quantization algorithm. All three of us used essentially the same one. The difference is where each of us decided to spend precision.

Three answers to the same question

Quantizing a 27B to 4 bits isn’t one decision, it’s a few hundred. Every tensor can stay high precision or drop to FP4, and each choice trades accuracy against bytes. Here’s what the three of us chose:

• NVIDIA put the MLP and lm_head in NVFP4 and kept self-attention and Gated-DeltaNet at FP8 — about 22% of the linear weight at double precision.
• QUASAR put everything in NVFP4 except lm_head, which they kept at BF16, and recovered accuracy with one epoch of quantization-aware distillation against the frozen BF16 teacher.
• We quantized everything, including lm_head, keeping only the embeddings and a few tiny Gated-DeltaNet gates at higher precision.

Nobody used flat NVFP4. Per-tensor allocation is the method — the scalar quantizer is essentially a solved commodity.

The lm_head decision, and why it paid twice

lm_head on this model is a 248,320 × 5,120 matrix — 1.27B parameters. In BF16 that’s 2.54 GB; in NVFP4 it’s 0.71 GB.

It is also read on every single token. And decode on a 5090 is bandwidth-bound: we measured our decode at ~1.55 TB/s against the card’s 1.79 TB/s spec, about 87% of peak. When you’re that close to the memory ceiling, bytes are time.

So quantizing lm_head bought two things at once. It made decode faster, because there are fewer bytes to stream per token. And it freed 1.8 GB, which becomes KV cache — context.

That single choice is most of the gap in the table. QUASAR’s BF16 lm_head is why their checkpoint cannot serve this model’s native 262,144-token window on a 5090 — vLLM refuses and caps them at 257,152. Ours serves the full window with 305,834 KV tokens, and retrieves a needle correctly at 256,130 tokens while still decoding at 60.8 tok/s there.

Why NVIDIA’s build is slower on a 5090, and why they’re not wrong

NVIDIA’s card says the target is GB300. That’s sm_100 — Blackwell with the tcgen05 tensor core path. The RTX 5090 is sm_120, which uses mma.sync.aligned instead: no tensor memory, no multicast, TN layout only.

Those two chips price precision differently, and I measured how differently. On our actual GEMM shapes, leaving NVFP4 for the next instruction tier costs 1.5–2.1×:

That’s a cliff, not a gradient. NVFP4 and MXFP4 run on mma.sync.aligned.kind::mxf4nvf4.block_scale at 2× Ada-FP8 rate; MXFP6 and MXFP8 share kind::mxf8f6f4.block_scale at 1×. There is no intermediate step. So on sm_120, NVIDIA’s FP8 attention and GDN costs about 18% of decode and 16% of prefill — and in the end-to-end measurement it came out at 25%, because their checkpoint is also 1.17× larger overall and decode tracks bytes almost linearly.

They didn’t make a mistake. They optimized for their hardware. The allocation simply doesn’t transfer to a consumer card, and that’s the whole opportunity.

The drafter: where 420 tok/s on code comes from

Quantization gets you to 92.9 tok/s. The next multiplier is speculative decoding, and it is the part that can’t be downloaded separately — because the drafter has to be trained against this checkpoint.

DSpark is a 5-layer semi-autoregressive block drafter with block size 7: it proposes up to eight tokens at once, and the target model verifies all of them in a single forward pass. Verification is exact, so speculation changes speed and not outputs.

We went through four versions, measured on the same 240 held-out prompts:

+31.7% over the model’s own built-in MTP head, using a quarter of the memory. That result is why we deleted the MTP head from the weights entirely — once the drafter beats it decisively, MTP is 0.85 GB of dead bandwidth.

Two things produced the gain.

Quantizing the drafter itself. The drafter runs on every speculation step, so it’s on the same bandwidth-bound critical path as the target. NVFP4 took it from 2.72 GB to 1.41 GB, and decode from 141.99 to 150.74 tok/s — +6.2% from shrinking a model that isn’t even the one answering.

Fixing a train/serve mismatch. This is the one worth stealing. The v1 drafter was trained on a conventional corpus. But the model in production sees something quite specific: XML tool calls with a real tools array, think blocks, a reasoning-effort preamble. The drafter had never seen the format it was asked to predict.

So we re-rendered the training corpus exactly as the model is served — same template, same tool syntax, same preamble — and retrained. Acceptance went from 2.761 to 2.904 overall. But the interesting part is where it landed:

A 5.2% overall gain that is +17.9% on agentic traffic. The general-purpose average hid almost all of the value. If your model serves agents, evaluate your drafter on agent traces or you will under-measure it by a factor of three.

The per-workload spread explains the headline number too:

Code accepts 6.78 of a possible 8 tokens per block. Code is predictable — closing brackets, boilerplate, repeated identifiers — so the drafter is right almost every time. Open-ended chat accepts 2.32 and gets 1.80×. Speculation is not one speedup; it’s a distribution, and its shape follows how predictable your traffic is.

(These come from SparkInfer’s bench harness at short context, where the autoregressive baseline is 97.1 rather than the server’s 92.9. The ratios transfer; the absolute numbers are their own harness.)

Runtime co-optimization: the half you can’t download

Every competitor in the table above ships a checkpoint. We ship the checkpoint and the engine, and the compounding between them is most of the remaining advantage.

As of v0.5.6, SparkInfer decodes this checkpoint faster than vLLM at every concurrency from 2 to 32 streams:

Same box, same checkpoint, same workload, three reps per point — the reps agree within 0.4% for SparkInfer and 1.1% for vLLM, and vLLM is given its best memory setting rather than the default.

That is a reversal. Four months ago the same comparison had us flat at ~77 tok/s no matter how many clients arrived, while vLLM scaled to 7× that. Three things closed it.

The model had no batch dimension. The scheduler was real — admission, chunked prefill, a proper worker loop — but decode called into the model one sequence at a time, so N clients were time-sliced rather than batched and every token still paid a full pass over the weights. Giving decode a single packed forward was the structural fix.

The packed forward then had to reach the tensor cores. A batched path that falls back to a row-wise GEMV is batched in name only. Moving the wide projections, the Gated-DeltaNet projections and the LM head onto the block-scaled NVFP4 GEMM — and making the batch ramp to full width immediately instead of one or two rows per iteration — is what turned the structure into throughput.

The hybrid architecture quietly penalises generic engines. vLLM has to set its attention block size to 1,568 tokens on this model so that a single page can hold a Gated-DeltaNet recurrent state. That leaves room for far fewer concurrent sequences: at 32 streams it reports a maximum concurrency of 27 even at gpu_memory_utilization=0.94, so it’s memory-bound before it’s compute-bound. Fitting all 32 would need about 0.99, which a 32 GB card doesn’t have. An engine written for this model’s 16-paged / 48-recurrent split doesn’t pay that tax.

The same split produced the subtlest bug we’ve had. 47 of the 48 Gated-DeltaNet layers were taking their decay gate and update rate from layer 0’s normalisation — an FP4 norm deferral skipped the refresh those projections read. Everything ran; the numbers were simply built on one layer’s statistics. Fixed from 32,768 tokens up.

And a few more that only exist because of the hybrid:

Prefix caching had to restore recurrent state, not just KV. Reusing a cached prefix is standard — keep the KV blocks instead of freeing them. On this model that’s a quarter of the work: 48 of 64 layers carry state that decoding mutates in place, so keeping the KV alone leaves the model holding the previous request’s history. With both handled, a repeated 32,022-token prefix went from 5.75 s to 0.30 s — 19× — with a cache-miss control to prove it wasn’t warm-up.

A fused prefill kernel was silently wrong above 2,048 tokens. It wrote its intermediate plane with one row stride and read it back with another; the two agree only below the padding boundary, so everything short passed and everything long came back subtly wrong and non-reproducible.

A long-context optimisation was quietly lobotomising agents. A sparse-attention path engaged above 16,384 tokens and attended only the first block plus the last 4,096 tokens — fast, and it meant a long agent session couldn’t read back its own earlier tool results. A 44K prompt with an ERROR line 20K back answered “there is no ERROR line.” It’s now opt-in rather than default.

None of those are quantization problems. All of them decide whether the quantized checkpoint is usable, and none would be found by someone who only ships weights.

(The vLLM comparison is engine-level: both sides run in-process, not through HTTP, and the release explicitly doesn’t claim the HTTP servers, more than 32 streams, longer prompts, or any other engine at concurrency.)

What I got wrong

Three things, and they’re the useful part.

I thought learned FP4 scales were unclaimed. I proposed it as our novel contribution. Then I found NVFP4_FOUR_OVER_SIX_CFG in the ModelOpt version already installed on my own box — plus NVFP4_W4A4_WEIGHT_LOCAL_HESSIAN_CFG and an MSE scale sweep. Four Over Six does exactly what I’d proposed, adaptively choosing per block whether to scale to 4 or to 6. I had printed that config list hours earlier and not read it.

I proposed an MXFP6 fallback tier. Better accuracy than FP4, 24% smaller than FP8, same instruction. Then I checked QuTLASS and there is no FP6 kernel — only mxf4, nvf4, mxf8. CUTLASS’s sm_120 headers do carry e2m3/e3m2, so it’s buildable, but it would need writing. And nothing in our accuracy data says we need it.

I hypothesized long-context quantization drift. With 48 of 64 layers being Gated-DeltaNet, weight error should compound through the recurrence, so short-context benchmarks would systematically under-measure damage in hybrids. It’s a good story. I ran AA-LCR — 100 questions at ~95K tokens each — on all three checkpoints and got 55/100, 55/100, 55/100. Dead heat. Hypothesis dead, one day spent, worth it.

What I’d tell you if you’re doing this

Measure the cliff before designing around it. The sm_120 tier structure decided more than any algorithm choice. One day of microbenchmarks on the real GEMM shapes would have saved me two wrong proposals.

Your biggest tensor isn’t your most important one. The MLP is 73% of the weight, so it must stay in NVFP4 — that’s not a decision, it’s arithmetic. The interesting choices are in the 27% that’s cheap to promote.

Publish the negative results. The long-context drift hypothesis failing is genuinely useful: it means aggressive flat-NVFP4 is safe on hybrid architectures, which nobody had shown.

Verify against published competitors, not against your own baseline. I’d have been happy with our numbers indefinitely. Running NVIDIA’s checkpoint on the same box is what turned “we think this is good” into “here is the gap and here is why it exists.”

The honest caveats

Published BF16 references disagree across harnesses — NVIDIA reports 88.92 on GPQA Diamond, QUASAR reports 91.41, for the same weights. A 2.5-point spread from harness alone. So the numbers in this post compare only to each other.

AA-LCR here is scored with a deterministic salient-token proxy, not Artificial Analysis’s LLM judge. It ranks these three consistently; it is not comparable to their leaderboard.

At n=198, the GPQA interval is ±6 points. A real 2-point gap could hide inside it. And all three arms still carried 41–45 truncations.

Where this goes

The checkpoint is on Hugging Face, and the engine that serves it, SparkInfer, is open source.

The next real question isn’t a better quantizer. It’s whether you can allocate precision against measured kernel latency and decode bandwidth on one specific card, instead of against a bit budget — because every published allocator optimizes bits, and bits can’t see a 2× cliff.

That’s the part nobody has done yet. It’s also the part you can only do if you ship both the checkpoint and the runtime — which, as of v0.5.6, is what beating vLLM at every concurrency from 2 to 32 streams on this model actually took.