How we got Qwen3.8-27B to 92.9 tok/s and a 360K window on one RTX 5090

Our NVFP4 build of Qwen3.8-27B runs a 27-billion-parameter model on a single consumer GPU at 92.9 tok/s, holds a 360,000-token context on 32 GB, and reaches 420 tok/s on code with a matched speculative drafter.

This post is the how. Not the benchmark table — the reasoning that produced it.

Start by finding out what you’re actually limited by

Everything downstream follows from one measurement.

At batch size 1, this model is weight-bandwidth bound, not compute bound. An early build streamed 18.80 GiB of weights per token at 81.6 tok/s. That works out to a 1.65 TB/s read rate against the RTX 5090’s 1.79 TB/s spec — about 92% of peak memory bandwidth.

That single number reframes the whole problem. Decode reads the entire model once per token, and we were already near the memory ceiling. So:

Bytes removed convert almost linearly into tokens per second.

Not “quantization might help.” Quantization is the optimization, and every byte you leave in high precision is a direct tax on decode. There’s no clever kernel that gets around it — you’re waiting on DRAM.

This also tells you what won’t help. Fusing ops, better scheduling, smarter batching at bs=1 — all of it is rearranging work that isn’t the bottleneck.

The tensor that mattered: lm_head

lm_head on this model is a full-vocabulary GEMM: 248,320 × 5,120, about 1.27B parameters. In BF16 that’s 2.54 GB. In NVFP4 it’s 0.72 GB.

Crucially, it is evaluated on every single token. It was the last large BF16 block sitting on the per-token critical path.

So we quantized it. That one decision moved the checkpoint from 20.59 GB to 18.77 GB and decode from 81.6 to 88.45 tok/s — an 8.4% gain from a single tensor.

The instructive part is the tensor we didn’t quantize for speed. The embedding table is exactly the same size — 248,320 × 5,120 — but it’s a gather, not a GEMM. You read roughly 10 KB per token, not 2.54 GB. Quantizing embeddings saves disk and VRAM capacity; it does nothing for decode speed.

Same shape, same parameter count, completely different optimization value. Size tells you almost nothing; access pattern tells you everything.

Deleting weights that no longer do anything

The checkpoint shipped with a Multi-Token Prediction head — roughly 0.85 GB of tensors.

Once you’re using a proper speculative drafter, no supported serving path uses MTP. Our DSpark v2 drafter beats it by 31.7% while using a quarter of the memory. So the MTP head was 0.85 GB being streamed for nothing.

We removed it entirely. Final checkpoint: 17.92 GB in 2 shards, down from 20.59 GB in 3.

Deleting weights is an underrated optimization. It has no accuracy cost when the weights are genuinely dead, and on a bandwidth-bound workload it’s pure profit.

Why NVFP4 specifically, and not “4-bit”

NVFP4 isn’t just a small number format. On Blackwell it’s a native tensor-core operand: 4-bit E2M1 values with UE4M3 scales per group of 16.

That matters because the alternative — GGUF-style k-quants, or any INT4 scheme — has to dequantize before it can compute. On older hardware that unpacking can eat more than half of execution time; there’s a whole line of research devoted to hiding it.

On the RTX 5090 (sm_120), mma.sync.aligned.kind::mxf4nvf4.block_scale consumes packed FP4 directly. There is no unpack step to hide, because there is no unpack step.

We measured what that’s worth on our actual GEMM shapes. Against BF16 at prefill:

BF16 saturates around 220 TFLOP/s on this card. NVFP4 reaches ~1,200.

And there’s a cliff, not a gradient: NVFP4 and MXFP4 run on the 2× instruction tier, while MXFP6 and MXFP8 share a 1× tier. Dropping a tensor out of NVFP4 costs 1.5–2.1× on that GEMM. So “keep this one at higher precision for safety” is an expensive instinct on this hardware.

Context: the other half of the memory

Qwen3.8-27B is a hybrid — 64 layers, of which only 16 are full-attention and 48 are Gated-DeltaNet. Only those 16 carry paged KV.

That works out to 16 × 2 × 4 heads × 256 dim = 64 KiB per token at BF16. The 48 linear-attention layers carry a fixed-size recurrent state instead, which scales with concurrency rather than sequence length.

So context capacity is a straight subtraction: whatever the weights don’t occupy becomes KV. Every gigabyte saved on weights buys roughly 16,000 more tokens of BF16 context, or double that at FP8.

That’s why the lm_head decision paid twice — 1.8 GB freed is both faster decode and more context. We serve the full native 262,144-token window with 305,834 KV tokens available, and hold 360,000 at the ceiling. A needle placed at 256,130 tokens comes back correctly, with decode still at 60.8 tok/s that deep.

Speculation, matched to this checkpoint

The last multiplier isn’t quantization at all.

A DSpark v2 drafter — a 5-layer semi-autoregressive block drafter with block size 7, trained against this exact NVFP4 checkpoint — gives 2.73× overall and 4.32× on code, peaking at 420 tok/s.

The training detail that mattered: the drafter was retrained on a corpus rendered exactly as the model is served — XML tool calls with a real tools array, think blocks, the reasoning-effort preamble present. Closing that train/serve mismatch moved acceptance from 2.761 to 2.904 overall and +17.9% on agentic tool calling.

Speculation is lossless by construction: the target verifies every drafted token, so it changes speed and not outputs.

The environment

Everything above was built and measured in a specific setup, and the details matter for reproduction.

Hardware. NVIDIA GeForce RTX 5090, 32,607 MiB, compute capability 12.0 (sm_120a), driver 595.84. Host: 24 cores, 62 GB RAM. These were rented, ephemeral boxes — assume the machine disappears, and back up anything you write on it immediately. (We lost a benchmark harness that way once.)

sm_120 is worth stating explicitly because it is not the same as datacenter Blackwell. The B200/GB300 are sm_100 with the tcgen05 tensor-memory path; the 5090 uses mma.sync.aligned, with no tensor memory, no multicast, TN layout only, and four supported tile shapes. Kernel results from one do not transfer to the other.

Quantization. NVIDIA Model Optimizer at git c4129b6, NVFP4 W4A4 with group size 16, FP8 KV cache via fp8_cast, calibrated on 128 image-text samples.

Left in BF16 deliberately: the vision tower (333 tensors, all preserved so image and video need no re-export), embeddings, and the Gated-DeltaNet conv1d / in_proj_a / in_proj_b.

One trap worth repeating, because it cost us an afternoon: excluded modules must be listed in two places — hf_quant_config.json (exclude_modules) and config.json (quantization_config.ignore). Miss the second and you get Parameter lm_head.input_scale not found at load time, which looks nothing like the actual problem.

Serving. SparkInfer, our own C++/CUDA engine, built with -DCMAKE_CUDA_ARCHITECTURES=120. Also runs unmodified on SGLang and vLLM. The published Docker image is ~1 GB — built in a CUDA devel image and shipped on CUDA base, which is what took it from 4.14 GB to 1.02 GB.

Verification. Every number here was measured client-side over the streaming API, never from engine telemetry. That’s not pedantry: our server’s own decode_tps field divides by a total that includes time-to-first-token, so at 32K context it reports 26 tok/s when true decode is 94. If you benchmark from the server’s own counters, you will publish wrong numbers.

What I’d take away

Measure the ceiling first. 92% of memory bandwidth told us quantization was the only lever that mattered. Without that number we’d have spent weeks on kernels that couldn’t have helped.

Access pattern beats size. lm_head and the embedding table are the same shape. One was worth 8.4% of decode; the other was worth nothing.

Delete before you optimize. The MTP head was 0.85 GB of weights doing no work.

Know your exact silicon. sm_120 and sm_100 are both Blackwell and they price precision differently. An allocation tuned for a B200 gives away roughly a quarter of your decode on a 5090.

The checkpoint is on Hugging Face and the engine is open source.