2026-09-24

How I forked vLLM and made Qwen3.8-Flash-Next decode 1.6x faster on one GPU

I run Qwen3.8-Flash-Next on one RTX PRO 6000 Blackwell as my daily local model. Over two days (Sep 23 to 24, 2026) I tuned the serving stack. Decode went from about 195 to 311 tok/s on short chat answers, step time from 19.0 ms to about 11.5 ms, and prefill from 7.9K to 12.9K tok/s at 100K context. The n-gram table stays in swap the whole time.

Below is each change, why I made it, and what it measured. Everything is in the fork: TheRealBluesun/vllm-flash-next, branch flash-next-tuning. The full report is flash-next/REPORT.md.

The model and the constraint

Flash-Next has an unusual architecture. From the report:

• 48 layers: 36 GDN linear-attention, 12 QSA sparse-attention.
• Hyper-connections with 4 residual streams.
• MoE, 512 experts, top-10. Experts are NVFP4 in the checkpoint; dense projections are BF16.
• A per-layer n-gram embedding table (PLE) on layer 1: about 320M rows of 160 bytes, about 48 GB in FP8.
• A native MTP draft layer, used for speculative decoding.

The card: RTX PRO 6000 Blackwell, sm_120, 96 GB, about 1.79 TB/s, 188 SMs, 128 MB L2, no tcgen05/TMEM.

My constraint: the PLE table stays in pageable memory and swap, unpinned, so the workstation stays usable. That rules out upstream vLLM's pinned-memory UVA offload (vllm#54371). The base is peakcrosser7/vllm, branch release/qwen38next_offload at 357e054, which serves the table from a CPU worker process. My 24 commits sit on top of that.

Goal: make decode and prefill as fast as possible on this card without a noticeable quality loss.

How I worked

Profile with nsys, find where the step time goes, change one thing, measure. Every change sits behind an env var registered in vllm/envs.py and does nothing unless set, so removing the tuning drop-in restores stock behavior. I measure speed as ms per decode step, because tok/s depends on how many draft tokens get accepted on a given text. I check quality as perplexity against the BF16-dense model on 40 x 512-token passages and 300 x 60-token passages.

Results

Same benchmark (tools/bench.py), median of 3 runs.

Quality cost: +0.4 to 0.65% perplexity on long passages and about +0.3% on short ones, all from the FP8 dense weights plus about 0.1% from the FP8 target lm_head. Every other change is bit-exact or only touches draft tokens, which the target verifies.

1. Making swap fast enough

Commit 74232c6. With the table partly paged out, each small index_select gather takes one serial major fault per cold row. That stalled the GPU about 4 ms per decode step.

• Swap-aware prefetch (VLLM_PLE_PREFETCH, vllm/v1/ple_offload/prefetch.py). Before a gather, the worker issues MADV_WILLNEED for every page it will touch via process_madvise, 1024 iovecs per call, so swap-ins run in parallel instead of one at a time. PLE stall went from about 4.3 to about 1 ms per step.
• Prefill hints (VLLM_PLE_PREFILL_HINT, connector.py, worker.py). When a prompt is admitted, the rest of it goes to the worker, which prefetches its rows in a background thread while earlier chunks compute.
• NumPy n-gram ids (VLLM_PLE_NP_IDS). 2x faster than torch, bit-identical, fuzz-tested.
• Combined effect at commit time: decode step 19.8 to 18.0 ms, 16K prefill 2.06 to 1.52 s.
• vm.page-cluster=0 so a swap-in reads 1 page instead of 8 for random 160-byte rows.

Later, commit 21b722a added pure-Python ids for decode batches of 32 tokens or fewer (VLLM_PLE_PY_IDS), computing only the context positions each token needs. Ids went from 0.16 to 0.07 ms per request.

This doesn't need fast storage. The gain comes from turning serial faults into parallel reads.

2. FP8 for the dense layers

The checkpoint ships dense projections in BF16. Online FP8 (weight-only, 128x128 blocks) for those layers was the biggest single decode win. Tiny layers, the router, shared experts, the indexer, the HC inject and the vision tower stay BF16. This is the only change that affects quality, about +0.5% perplexity.

Commit 6053717 extends it to the QSA fused qkv projection (VLLM_QSA_QKV_ONLINE_QUANT) for a small decode gain.

3. Cheaper drafting

Speculative decoding uses the model's own MTP layer. I went from k=3 to k=4 (+7 to 9% on chat).

Commit 9753db9. A greedy drafter only needs an argmax, and the target verifies every draft. So:

• FP8 draft lm_head (VLLM_DRAFT_HEAD_FP8, fp8_draft_head.py), a per-row FP8 copy run as a Triton GEMV.
• Hot-token draft vocab (VLLM_DRAFT_VOCAB_FILE). Drafting is restricted to the 98K most-used tokens, which cover 99.6% of chat. 64K and 70K lost too much acceptance.
• Draft phase went from about 2.9 to 2.1 ms per step with acceptance flat. Neither change can alter outputs. The vocab subset is English-centric.
• FP8 target lm_head for decode-sized batches (VLLM_LM_HEAD_FP8): -0.5 ms per step, about +0.1% perplexity.

Commit 961a7af. The draft layer's quant config is rebuilt from the draft model config, which never carried --quantization-config, so the MTP layer was silently staying BF16. VLLM_DRAFT_ONLINE_QUANT attaches FP8 dense plus FP8 128x128-block experts to the draft only. -0.3 to 0.4 ms per step, about 3%.

4. Narrow GEMMs

Commit 30616f6. For BF16 layers with 256 or fewer outputs at decode batch sizes, cuBLAS picks 8-CTA WMMA kernels, for example 13.8 us for a 96x2560 layer on sm_120. VLLM_NARROW_GEMM routes them through a split-K Triton kernel at about 2.8 us (narrow_gemm.py).

5. Hand-written kernels

At M around 5, decode linears do nothing but stream weights. I wanted DRAM busy across kernel boundaries.

PDL weight-prefetch GEMV (VLLM_PDL_GEMV, commit c4e1511, pdl_gemv.py). A split-K GEMV for FP8 128x128-block W8A16 and BF16. Each CTA loads its whole weight tile before griddepcontrol.wait and triggers dependents right away, so with programmatic dependent launch the next kernel streams its weights while the current one finishes. Partials reduce with fp32 atomics. The last CTA per tile writes the output and re-zeroes a fixed-size workspace, so CUDA graphs never see it replaced.

• The first version was slower (0.84x) because it gathered one fp32 scale per weight element. Switching to one broadcast scale per block fixed it.
• In a prototype, a 96-GEMV chain went from 64% (Marlin + cuBLAS) to 78% of DRAM roofline.
• In the server it only paid off once the neighbouring kernels were PDL-friendly. The HC Triton kernels now trigger right after their wait, and the GDN decode post-conv CUDA kernel got an early launch_dependents through a runtime-built copy (pdl_ext/gdn_post_conv_pdl.cu, bit-identical output and state).
• Corpus 12.39 to 11.88 ms per step, chat 12.39 to 11.60. Costs 2.7 GB of KV cache for the row-major FP8 copy, still above max context.

L2-resident draft layer (VLLM_L2_DRAFT, commit d59cd5d). The MTP draft layer's ~118 MB of dense weights get read 4 times per step. CUDA graphs don't capture the persistence API, so this uses per-load eviction hints. Simulation showed evict_last on the draft weights protects nothing by itself. Everything else has to stream evict_first; even 31 MB of default-policy traffic between passes undoes it. About -0.07 ms per step.

Fused MoE routing (VLLM_FUSED_ROUTE, commits a71d447 and 36e1a28, pdl_ext/moe_route.cu). For batches of 64 or fewer, one 1024-thread CTA replaces topkGating, moe_align_block_size, count_and_sort_expert_tokens and a fill. It does warp-per-token softmax, top-k by redux.sync arg-max on float-bit keys (lowest id wins ties, like the reference), then counts, padded scan, expert ids and scatter.

• Top-k cost 3.2 us until per-lane contiguous experts, vector loads and redux cut it to about 1.9 us.
• 7.4 to 3.7 us per MoE layer, -0.2 ms per step. Routing matches the reference exactly.
• torch.compile can trace the router, and an is_compiling() guard was silently selecting the reference path inside the compiled graph. The fix runs top-k as a registered custom op that decides at execution time. nsys confirms 52.6 moe_route_kernel calls per step and 0 reference routing kernels.

6. Config

• Prefill chunk (max_num_batched_tokens) 2048 to 8192: +7% prefill.
• MTP k=4.
• Online FP8 dense, as above.

What didn't work

Where the time goes now

About 12.6 ms per step under nsys:

• PDL GEMVs 3.5 ms, about 80% of roofline.
• Marlin MoE 2.9 ms.
• GPU idle about 1.9 ms. Roughly half is one gap per step, the PLE stall while the graph waits on the CPU worker's rows. Worker time per request is about 0.33 ms; the rest is hand-off latency (D2H, connector thread, zmq, worker, H2D).
• FP8 heads 1.1 ms, about 87% of roofline.
• GDN recurrence 0.7 ms, MoE routing about 0.3 ms, hc_combine_norm 0.45 ms.

Remaining levers, each about 1 to 3% unless noted: cut the PLE hand-offs by having the worker poll a GPU-written host flag; fold hc_combine_norm into its predecessor; then turn the TMA GEMV back on; a calibrated NVFP4 dense checkpoint; and a persistent decoder-layer megakernel, the only big jump left (up to 30 to 40%) and a large project.

Gotchas

• vLLM's torch.compile cache is keyed on config plus registered env vars. Unregistered toggles reuse a stale graph (commit 94dfbcd).
• Python in the router can be traced. Make runtime decisions inside a custom op.
• INFO logs during CUDA-graph capture don't appear. Confirm kernel paths in nsys.
• sm_120 + TMA: cp.async.bulk to a .shared::cluster destination compiles with a software fallback that makes the driver reserve about 3.6 GB of VRAM for stack. Use .shared::cta.
• PDL: a kernel only overlaps a predecessor that triggers early, and any non-PDL kernel is a barrier.
• With a second GPU in the box, pin the serving card. Otherwise FlashInfer JITs for both architectures at startup and runs the host out of memory.
• Yama ptrace_scope=1 blocks the pidfd_getfd CUDA IPC uses to import the PLE worker's buffers; opt in with PR_SET_PTRACER (commit 1f0b645).

Upstream

I posted the swap-prefetch findings on vllm#54070 (disk-backed PLE tables). The /start_profile crash I fixed in cdc4f07 was fixed the same way upstream in vllm#55237. If the CPU-worker offload design lands upstream, these could go with it: PLE swap prefetch, the FP8 draft head with a vocab subset, online quant for the MTP draft layer, the fused routing kernel (generic for softmax top-k at small M), and the narrow split-K GEMM.

Code

github.com/TheRealBluesun/vllm-flash-next, branch flash-next-tuning. flash-next/ has the report, roadmap, lab notes, deploy config, benchmark and profiling tools. Opus 5.5 was a big help implementing the kernels.

Code: TheRealBluesun/vllm-flash-next