How I made Qwen3.8 27B decode 1.8x faster at 109K context on one RTX 3090 Ti
I run Qwen3.8 27B (4-bit, W4A16) on two 24 GB Ampere cards, an RTX 3090 Ti and an RTX 3090, each with at least 128K of context. Over three days (Sep 24 to 26, 2026) I took the stack apart. At 109K tokens of context, decode went from 52 to 94 tok/s, and a cold 109K prefill went from 277 s to 113.6 s.
Each change is listed below with the reason for it and its measured effect. The code is in TheRealBluesun/HyperQwen-3090ti, branch qwen27-8001, directory local-3090ti/. Every measurement is in PROGRESS.md there.
Starting point and constraint
The base is HyperQwen (syv-ai), a patch series on vLLM 0.29.0 for serving big Qwen models on consumer GPUs. My patches (01 to 18) apply on top of it (patches/apply.sh). The model is syvai/Qwen3.8-27B-DFlash2-W4A16, with DFlash2 speculative decoding (a drafter plus an n-gram style lookup) and an int8 KV cache.
The model is a hybrid. Most layers are GDN linear attention with a recurrent state, and the rest are full attention. The drafter has its own sliding-window attention layers.
The constraint is lossless output. Every change either produces bit-identical output or is checked against a non-speculative reference. No lower-bit weights.
How I measured
tools/bench27.py: streamed, engine-neutral, greedy, median of 3 runs.LONG=1adds a ~109K-token test.- ms per verify step, read from vLLM's draft counters. It is stable to about ±0.2 ms. tok/s is not, because it moves with acceptance.
- nsys for kernel time. For correctness: greedy output comparisons, valid JSON on a replayed workload, needle retrieval at 100K, GSM8K, and an oracle check that teacher-forces outputs through a non-speculative bf16 model (
lossless.py).
Results (3090 Ti, as cloned vs patches 01 to 18)
| Metric | As cloned | Patched | Change |
|---|---|---|---|
| Decode, code / prose / explain / json (tok/s) | 230 / 106 / 163 / 264 | 266 / 132 / 193 / 305 | 1.16 to 1.25x |
| Prefill, 18K prompt (tok/s) | 956 | 1,313 | 1.37x |
| Prefill, 109K prompt | 277 s (395 tok/s) | 113.6 s (962 tok/s) | 2.44x |
| Decode after 109K of context (tok/s) | 52 | 94 | 1.81x |
The baseline ran at a 300 W power cap and the final numbers at 350 W, so part of the short-context gain comes from power (see below). For agent workloads the prefix-cache patches matter just as much: a turn re-prefills about 26 already-seen tokens instead of about 394.
0. Power first
Before writing kernels I checked why the Marlin GEMMs ran slower in the server than in isolation. During sustained decode the card sat at its 300 W cap, with the SM clock around 1,500 MHz instead of ~2,000. A gate_up GEMM went from 107 us in a burst to 140 us sustained. Decode was power-limited, and fusing kernels wouldn't have fixed that.
Power-limit sweep on the 3090 Ti:
| Limit | ms/step | Decode vs 300 W |
|---|---|---|
| 300 W | 24.32 | |
| 325 W | 22.47 | +8.2% |
| 350 W | 22.00 | +10.5% |
| 400 W | 21.95 | +10.8% |
The knee is 325 to 350 W. I set both cards to 350 W.
1. Verify attention in CUDA (patch 06)
At long context, the verify step's attention over the int8 KV cache dominates decode time. Nsight showed the Triton kernel at 1 CTA per SM, 254 registers, 1.56 waves and 18% DRAM utilization.
Patches 01 and 02 first re-tuned the Triton configs for sm86 (verify: 32 splits, 64-token tiles; prefill: 128-row blocks). That took decode at 109K from 52 to 66 tok/s and 109K prefill from 277 to 162 s.
Then I wrote a CUDA split-KV verify-attention kernel that reads the int8 per-token-head cache directly (qwen27_sda.cu, JIT-built). Iterations, per layer at 110K context:
| Kernel | 110K | 64K |
|---|---|---|
| Triton at the start | 1,176 us | 710 us |
| v1/v3: cp.async double buffering, Q in registers, 2 CTAs/SM | 552 us | 333 us |
v4: fp16 Q.K^T, cp.async race fixed (commit 806c17a) |
536 us | 325 us |
v6: softmax in registers (commit 7bd90c7) |
508 us | 306 us |
v7: two-level int8 Q, s8 MMA for Q.K^T (commit 705e86c) |
355 us | 221 us |
v7 is 3.3x faster than the Triton kernel at 110K. Over that night, 121K decode went from 61 to 94 tok/s. Error against an fp32 reference didn't change (~2.4e-3).
2. Prefill attention through the same kernel (patch 12)
Commit c90d622. Run in direct mode (one segment per query tile, normalized bf16 output), the v7 kernel replaces vLLM's Triton attention for all non-verify attention on the int8 cache. It reaches 64 to 67 TFLOP/s effective, against 31 for Triton.
| Context | TTFT before | TTFT after |
|---|---|---|
| 18K | 15.5 s | 14.3 s |
| 34K | 33.3 s | 28.9 s |
| 63K | 76.9 s | 60.7 s |
| 121K | 199 s | 139 s |
GSM8K (200 questions): 95.5%, vs 96.0% before. Prefill is now GEMM-bound.
3. Marlin tuned per decode shape (patch 08)
Commits 573a102, 2423b15, 22911da; the pipeline is in marlin_tune/. At M=8, stock Marlin already reaches 90 to 95% of the card's read ceiling on the big matrices, but only ~70% on the 16 MB o_proj. The patch adds a table of tile config (thread_k, thread_n) and CTA count per decode shape, one for M ≤ 8 and a second for 8 < M ≤ 16, enabled with VLLM_MARLIN_TUNE=1. Output is within 1 bf16 ulp of stock. Step time went from 24.13 to 23.54 ms, and weight GEMMs now run at about 92% of the memory-read floor.
4. Smaller kernel work
- Patch 03 (commit
7e5c792): small top-k path in the V2 sampler. Triton top-k/top-p cost 813 us per step, so I replaced it with a sort-freetorch.topk. 26.21 to 25.56 ms/step. - Patch 04 (same commit): split-KV attention for the drafter's sliding-window layers, down from 170 to 180 us to 58 us each. It also fixes the split-KV buffer, which was sized for a 256-request fallback. That frees 1.9 GB per card.
- Patch 07 (commit
b2edfb1): GDN attention metadata is built once per step instead of once per group. - Patch 09 (commit
1199731): small-batchsilu_and_mulin Triton, 4.9 to 1.8 us, bit-identical. - Patch 10 (commit
a332c5e): GDN spec decode reads q/k/v in place. 4 fewer kernels per layer, bit-identical.
Together, patches 08 to 10 took step time on the replayed workload from 24.13 to 23.06 ms.
5. Lazy GDN state commit (patch 14)
Commit b681d56, VLLM_QWEN27_LAZY_GDN=1. A spec-decode step of M tokens stores the recurrent state after every token (M x 1.57 MB per GDN layer), so the next step can resume from whichever token is accepted. The new CUDA kernel stores only the starting state plus a 133 KB log of the step's tokens, and the next step replays the accepted ones. It rounds through fp16 at exactly the points where the full path stores and reloads, so every state is bit-identical.
About 91% of decode rows take the lazy path. The recurrence went from 21.2 us (Triton) to 15.8 us per GDN layer, which is about +0.8 to 1% decode. That's less than the isolated benchmark suggested, because the state writes already overlapped other work.
6. Bug fix: corrupt output when the verify block shrinks (patch 13)
Commit 865bd6a. With more than 7 draft tokens, or with synchronous scheduling, greedy JSON output came out corrupt. The GDN conv kernel compared the previous step's num_accepted against the current block length, so any shorter block zeroed the conv output in every GDN layer. The fix bounds the check by conv-state capacity. On the oracle check with k=11 pinned, bad tokens dropped from 264 to 1 (the grammar forcing {). Production (async, k=7) wasn't affected.
7. Lookup drafting scan (patch 15)
Commit c167d18. The DFlash2 lookup's suffix scan ran as a single Triton program walking the whole history, costing 265 us per step at ~100K. I split it over 64 programs per request and folded the results with an int64 atomic max on the same packed score, so the proposal is exactly the old one. In isolation it went from 124 to 6 us at 105K. End to end it's worth about 1% at 100K.
8. Prefix caching for agents (patches 16 and 17)
I profiled a real agent harness on 12 coding and ops tasks. 37% of the prompt tokens it actually prefilled had been seen before.
Patch 16 (commit 60f247a). Every prefix-cache hit came up one full 864-token block short. With an EAGLE drafter, the KV at position p mixes in token p+1, so vLLM drops the last block. DFlash's context K/V at p depends only on the target hidden state at p, so nothing needs to be dropped. The patch adds a separate eagle_prefix_drop flag, false for DFlash. Re-prefilled seen tokens fell from 45% to 13% of new tokens, and TTFT went from 2.81 to 2.14 s (-24%).
Patch 17 (commit c708e59). The chat template drops the last turn's thinking, so each agent turn's prompt diverges from the previous one a token before its end. With block-granular hits, each turn still re-prefilled 394 tokens on average. vLLM supports fine-grained hits (--prefix-match-unit), but disables them when any KV group can't do fine-grained lookup. Here that group is the drafter's sliding window, and patch 17 adds the lookup for it. With --prefix-match-unit 32, re-prefill per turn dropped from 394 to 26 tokens, total re-prefilled from 93.5K to 5.6K (-94%), and median follow-up TTFT from 1.21 to 1.05 s.
What didn't work
| Idea | Result |
|---|---|
| Persistent megakernel (GDN + MLP layers) | Correct, but ~18% slower than the kernel sequence. Grid-wide barriers and static work split. |
| Custom W4A16 GEMV for M ≤ 16 | Ties Marlin at best |
| Tree verification / 16-token verify | Ceiling +7 to 12% tokens per step; a 16-token step costs +49% at 32K |
| Drafting from other requests' output | Offline ceiling +1.6%; real policies -0.1 to -1.9% |
| Fine-tuning the drafter | Qwen's own MTP head has the same per-position acceptance as DFlash2; the limit is the model's predictability |
| INT8 activations (W4A8) | TTFT -47%, but decode +5.8% ms/step and +4.1% perplexity upstream |
| int8 Q.K^T in the verify kernel (early) | +15% only, 16x the error with outlier q channels |
| Verify attention past v7 | Four variants, including warp-specialized; none beat it |
| Smaller drafter vocab (40,960 tokens) | -0.76 ms/step but fewer accepted tokens; -2.4% net |
| Fused CUDA GDN decode with bf16 state | -0.7% ms/step, not worth the quality risk |
| "Low" reasoning effort | The model thinks more, not less |
Where it stands
At 100K context a step takes 28.7 ms. Weight GEMMs account for 19.3 ms, at ~92% of the memory-read floor, and verify attention for 5.5 ms. Most of the rest is launch overhead, about 2.5 to 3 us across ~700 kernels per step, which only a persistent kernel can remove. Single-card lossless decode is close to its practical ceiling. Getting lower needs fewer weight bytes, and that means lossy weights.
Next
- Pipeline parallel over both cards (patches 19 and 20, Sep 29). Stock vLLM refused to run DFlash under PP, and patch 19 fixes that. 109K cold prefill drops from 113.6 s to 62.0 s, and decode after 109K goes from 94 to 101.8 tok/s. Short-context decode is still a few percent slower than on one card.
- A per-card persistent decode megakernel, redesigned without grid-wide barriers. Realistically that's worth about 1.2 to 1.3x single-stream decode, since the weights still have to cross the memory bus. The design note is in
ROADMAP.md. - An AMD Radeon R9700, to see how much of this transfers.
Code
github.com/TheRealBluesun/HyperQwen-3090ti, branch qwen27-8001. local-3090ti/ has the patches, the Marlin tuning pipeline, the benchmarks and the full progress log. Ultimately it was a lot of profiling with nsys and reading the data. Opus 5.5 was a massive help throughout; it's the first model I've used that competently implemented GPU code without stepping all over itself.