Running 125B MoE on Three 3090s - vLLM Optimization Edition

In the previous article, I got Qwen3.8-Flash-Next (a 125B-A6B MoE plus a 51B n-gram table) running in vLLM on three RTX 3090s with a 262,144 context and 80 tok/s single-stream decode. This is the sequel: the ten days between putting it into service and reaching the numbers on the current model card.

The results first.

* 80.0 was measured with prefix caching off. As I explain below, the production configuration with prefix caching on only did 59-73 tok/s.

The 3x3090+MTP column was measured with the release settings (--kv-cache-memory-bytes 550000000, 262,144 x 2). The 2-stream value used --max-num-seqs 2 (the release setting); the others used the earlier --max-num-seqs 4. The table in section 7 is a separate run on the development configuration (--kv-cache-memory-bytes 783000000, 2.85x), so its values differ slightly.

On top of the three patches from last time (vllm.patch, decode-01, decode-02), the release now has 11 more.

There were also two configuration changes that are not patches, and three approaches that were implemented but not adopted. They are described in chronological order below.

1. The first problem: 100 seconds to the first token

I put last article’s configuration into production and connected qwen-code. The first token of the first request took more than a minute.

The server log said:

Avg prompt throughput: 4206.0 tokens/s

Prefill looks fast. But reading the log in time order, prefill started at 08:43:44 and the first token came at 08:45:23: 100 seconds for 42k tokens.

Avg prompt throughput is not TTFT. vLLM v1 adds a prompt’s token count to the 10-second reporting window in which its first token appears. The value is just “prompt length / 10 s”. TTFT has to be measured from the client.

Another hint: later turns of the same conversation (55k-73k tokens) took 10-25 seconds. Only the first one was extremely slow.

Cause: page cache misses on the PLE table

As described last time, the 95.37 GiB PLE n-gram table is a file on NVMe, mmapped, and rows are gathered with np.take. Each token needs 16 rows (bigram and trigram, 8 heads each), 320 B per row, and every row is on a different page.

np.take reads the rows in order on one thread, so each page not in the page cache is a synchronous read at queue depth 1, one page at a time. Readahead is off, because last time I set MADV_RANDOM (without it NVMe reads are 23x larger).

Cold, a 512-token chunk takes 1.9 s, so prefill drops to 270 tok/s. If every page of a 42k-token prompt misses the cache, that is 82 chunks x 1.9 s ≈ 156 s. This is an upper bound; the measured 100 s is shorter. Part of the table (the 2.5 GiB below) was already cached, and n-grams that repeat within the prompt read the same rows, so fewer pages are actually read from NVMe than the estimate assumes. I did not measure how much each contributes.

Why the table is not in the page cache. The machine has 125 GiB of RAM, but the QSA host K/V pool holds 12 layers x 3.96 GiB = 47.5 GiB pinned (page-locked, never evicted). A 95.4 GiB table cannot fit. fincore showed 2.5 GiB of the 95.4 GiB in the cache. Every new text is read from NVMe.

The 122k configuration without offload left room for the table in RAM, so this only surfaced after moving to the 262k configuration.

The fix: prefetch the whole prompt when it arrives

Which table rows are read depends only on the token sequence. The whole prompt is known when the request arrives, so the pages can be requested before prefill gets there.

1. A hook in Qwen4ExpModelState.add_request hands a new request’s prompt tokens to a prefetch thread.
2. The prefetch thread computes the same hash as the GPU’s Triton kernel _ple_ngram_ids_kernel in numpy and produces the [num_tokens, 16] row numbers. It matches the kernel in int64 overflow (wrap), the sign of the remainder (torch.remainder), and treating context before the most recent EOS as EOS; a test confirms it is bit-identical. A wrong value would only prefetch the wrong page, since the forward never uses it.
3. For every 1024 tokens it finds the unique pages, merges adjacent ones into ranges and calls process_madvise(pidfd, iovec[≤1024], MADV_WILLNEED). The kernel queues the reads asynchronously, so the drive’s queue fills up.
4. The gather (_lookup) is not changed at all. It finds its pages resident or already being read.
5. When a request finishes or is preempted, the rest of its prefetch is dropped.

Note: under podman’s default seccomp profile, process_madvise returns EPERM without CAP_SYS_PTRACE. The compose file adds cap_add: [SYS_PTRACE]. Granting CAP_SYS_PTRACE also allows ptrace of other processes in the container. Without it the server still starts, logs PLE page prefetch uses per-range madvise, and falls back to one madvise per range: the same NVMe throughput, with about 20x the CPU for the prefetch thread. VLLM_PLE_PREFETCH=0 turns the prefetch off entirely.

Same machine, two different unseen prompts of about 37k tokens (vLLM Python source). The table was still on the drive it was on before the move to the P310 in section 3.7, the same condition as the 100 s above.

Warm does not change because advising a page that is already cached costs about 1 µs. Nothing is added to the warm path.

When every page misses the cache (16 pages per token), the prefill limit is set by the drive’s random read rate (just under 40k pages/s): about 2,400 tok/s. Going beyond that requires the table in RAM. The 4,430 tok/s with prefetch in the table above exceeds this limit. The Python source used for the measurement has many common n-grams, and some pages were already cached before the run, so fewer than 16 pages per token were read from NVMe. The same reason explains why 44 s without prefetch is shorter than the 100 s (42k) at the start of this section.

Cold measurements need unseen text. Send the same prompt twice and the second run has warm pages and a warm Triton cache. Each side of an A/B needs its own unseen file.

2. Prefix caching cuts the context by 30%

Last article’s configuration used --no-enable-prefix-caching. An agent like qwen-code sends the whole, growing conversation every turn. Without prefix caching, a 70k conversation is prefilled again from the start every turn.

With it enabled, the capacity was:

prefix caching off:  GPU KV cache size: 1,067,300 tokens  (4.07x)
prefix caching on:   GPU KV cache size:   734,862 tokens  (2.80x)

Meanwhile the GPUs had about 500 MB free. Free VRAM was not the cause.

Capacity is set by the number of blocks, and a block’s cost is host RAM

In section 5.2 of the last article, with QSA offloaded, the attention block size is raised to 12,144 tokens so that its page matches the mamba (gated delta-net state) page. Capacity then follows this formula, calibrated at four points:

blocks         = floor(--kv-cache-memory-bytes / 3,207,168)
capacity (tok) = blocks / R x max-model-len

• 3,207,168 B is one gated delta-net state (ssm fp32 3,145,728 + conv bf16 61,440).
• R is the number of blocks one 262,144-token request takes: 42 with prefix caching off, 61 with it on.

The difference of 19 is the number of linear-attention KV groups. Prefix caching needs --mamba-cache-mode align, which raises each group’s state from 1 page (the live state) to 2 pages (live state + checkpoint). The total is 19 x 2 + 22 (QSA, cdiv(262,144, 12,144)) + 1 (PLE short-conv) = 61.

So prefix caching uses no extra VRAM at all. With the same 550 MB / 171 blocks, the blocks per request went from 42 to 61, and 4.07 requests became 2.80.

So the number of blocks has to increase. The cost of one block:

That is 1,405 tokens per MiB of VRAM and 15,460 tokens per GiB of RAM. In practice, --kv-cache-memory-bytes sets the host RAM usage. The pinned host RAM is about 93 times the value.

Stopping the power-of-two round-up

So I raised --kv-cache-memory-bytes. At 590 MB, host RAM use was 112 GiB with 13 GiB free.

The cause is PyTorch’s caching host allocator, from section 5.4 last time: it rounds pinned allocations up to a power of two. A 4.24 GiB pool per layer becomes 8 GiB, 96 GiB for 12 layers. Last time I avoided it by staying at 550 MB (3.96 GiB per layer), just under 4 GiB.

Reading the source of the bundled PyTorch 2.13 (CachingHostAllocator.h), the round-up has a threshold: add pinned_max_round_threshold_mb:2048 to PYTORCH_CUDA_ALLOC_CONF and allocations above 2 GiB are not rounded.

In the end I stopped relying on an environment variable and allocate the host pools directly: an anonymous mmap of exactly the right size, pinned with cudaHostRegister. The pools live until the process exits, so no allocator cache is needed. Before pinning, it checks free RAM and refuses to start with the sizes printed if there is not enough. The three PP ranks allocate at the same time, so a file lock serializes them; otherwise two ranks could both pass the check and then run out of memory together.

The values:

With four 247,823-token requests resident at once: needle 4/4, 0 preemptions. This gives 262,144 x 4 with prefix caching on.

3. Decode, second round: production was not at 80 tok/s

Last time I counted, from the safetensors headers, the bytes one decode step reads: 5.83 GB/token, which at the 3090’s 936 GB/s gives a 160 tok/s ceiling. With PP=3 and three requests in flight, the throughput ceiling is 480 tok/s. 80 tok/s was half of that, so there seemed to be room.

3.0 Measurement method

Each improvement at this stage is 0.3-2 ms, less than the variation between boots. So the measurement method was fixed first.

• Fixed prompt, fixed seed, one warm-up run and three measured runs. Random prompts cause PLE page cache misses and move even the median.
• Look at both the median inter-token latency (ITL) and the mean tok/s. The mean includes outliers such as PLE page cache misses and NVMe stalls. The mean is what the user feels.
• A/B within one build, switched by environment variables. The same configuration gives different numbers on a different boot. That is why every patch here can be turned off with one variable.
• torch profiler reports longer kernel times than actual inside a FULL cudagraph. CUPTI adds a few µs per kernel, so the error grows with the number of kernels. It reported 12.15 ms in total while the actual ITL was 10.66 ms. The real GPU time is measured with CUDA events around the graph replay.
• For the host timeline I injected a tracer into all processes with a .pth file and recorded monotonic_ns. Match ranks by completion time, not start time.

3.1 About 70 with prefix caching on, 60 above 8k

Measured this way, the production configuration (prefix caching on) had a single-stream median ITL of 13.2-13.7 ms, a mean of 59-73 tok/s. The 80 tok/s from last time was with prefix caching off. With it on, mamba_get_block_table_tensor in mamba align mode launches 7 eager kernels per KV group (13 per rank): about 95 more kernels per step per rank.

On top of that, past 2,048 tokens of context the step grows by 2 ms, to 60 tok/s. With QSA offload off (a 32k configuration), context 0 / 8k / 24k were flat at 12.5 ms, so those 2 ms are entirely the PCIe read of QSA’s host K/V.

This corrects a mistake in the last article. I wrote that reading 48 MiB per token at 80 tok/s is a sixth of PCIe bandwidth and “can even overlap with compute”. It did not overlap. The rows the indexer selects in a layer cannot be read before that layer’s indexer has run, so the read cannot overlap with other work. At 45 tok/s, host-side waiting was larger and this was not visible.

3.2 The bottleneck was serial processing on the host (decode-03)

The hyper-connection fusion from 3.3, on its own, gave −0.4 ms at context 0 and −2 ms at 8k. GPU work decreased, but short prompts barely improved. Short-prompt steps were limited by something other than the GPU, at about 12.5 ms.

A host-side trace showed the cause. On every PP rank but the first, Worker.execute_model starts with irecv_tensor_dict(). Its first step is recv_object, which receives the tensor metadata over gloo and blocks until the previous rank has launched its forward and called isend. Only then do input preparation and attention metadata construction (3.2-3.5 ms of Python on this model) start.

old:  rank0 [prep][launch]──→ rank1 [wait recv][prep 3ms][launch]──→ rank2 [wait recv][prep 3ms][launch]
new:  rank0 [prep][launch]──→ rank1 [prep 3ms][recv][launch]──→ rank2 [prep 3ms][recv][launch]
                                ↑ overlaps the previous rank's GPU time

GPU time per rank is about 3.5 ms and the host preparation is about as long, so the preparation never overlapped the previous rank’s GPU work and sat on the critical path every step. That is why making the GPU faster did not go below 12.5 ms.

The fix wraps the receive in a DeferredRecvIntermediateTensors and delays it until the model runner reads .tensors, right before the forward. The order of device operations on the rank (receive, forward, send) is unchanged. If nothing reads it, it still receives after execute_model, to match the previous rank’s send.

3.3 Small GEMVs at about 20% of bandwidth (decode-04 / 05)

With the host bottleneck removed, GPU time becomes the limit. Computing the effective bandwidth per kernel, the large projections were fine (gated delta-net qkvz 740 GB/s, QSA qkv 750, INT8 lm_head 900). Only the small ones were slow.

I also tried Marlin’s VLLM_MARLIN_USE_ATOMIC_ADD=1. Marlin ignores it on sm8x + BF16, so it had no effect.

Hyper-connection (decode-04)

One hyper-connection GatedResidual is six kernels: the inject GEMV, down, silu, up and the gate mix, about 32 µs. There are two per decoder layer, so 96 per token.

compressed-tensors INT8 pack-quantized puts four values in each int32 word, least significant byte first. So weight_packed.view(uint8) already is q[N, K] (value + 128). No repack and no extra memory. I wrote two Triton kernels that read it directly:

• _hc_down_inject_silu_kernel: the down (INT8) and inject (BF16) GEMVs in one kernel, with silu fused into the store
• _hc_up_gate_mix_kernel: the up (INT8) GEMV, with the gate mix across the 4 HC streams done in registers

They round to BF16 at the same points as the unfused path. One GatedResidual goes from 32 µs to 13.3 µs; at 96 per token that saves about 1.8 ms.

Batches over 4 tokens (prefill) go through a tl.dot W8A16 GEMM on the same weights, faster than Marlin at 512 tokens (down 102 vs 120 µs). But a 64x64 tile took 176 µs and slowed prefill by 4%. I used 32x64. Parameters chosen for decode cannot be reused for prefill as they are.

MoE block (decode-05)

At one token, the MoE block spent more time on the small kernels around the expert GEMMs than on the GEMMs: the router GEMV (12 µs), topk_softmax (7 µs on one CTA), moe_align_block_size (7 µs, two kernels), the shared expert gate (7 µs, four kernels), moe_sum and the final add. The expert weights themselves are 25 MB and take 42 µs.

• The router and the shared expert gate are one GEMV (512 rows + 1).
• The CTA that finishes last writes the top-k, the renormalization and the Marlin block alignment. Each CTA takes a ticket with atomic_add(sem="acq_rel") after tl.debug_barrier(), and only the last one does the post-processing. No separate kernel for grid synchronization.
• With norm_topk_prob=True, the full softmax denominator cancels out. A softmax over the 10 selected logits gives the same weights.
• For top-k, each BF16 logit becomes an order-preserving 16-bit integer, with 65535 − expert id packed below it into a 32-bit key. One max gives the value and the id at once, and ties go to the lower id (like vLLM’s original kernel). Faster than two argmax passes.
• The routed expert GEMMs call vLLM’s _fused_marlin_moe unchanged.
• The INT6 shared expert’s compressed-tensors layout (32 values tightly in six int32 words) is complex to unpack in a GEMV, so at load time it is rearranged into a 4-bit plane and a 2-bit plane, same number of bytes. One kernel does gate_up + SiluAndMul, one does down + the final combine (sum of routed + sigmoid(gate) x shared). 13.4 / 7.0 µs -> 7.8 / 5.3 µs.

3.4 PLE for a freshly sampled token (decode-06)

The prefetch in section 1 only covers the prompt. The n-gram rows of a token that was just sampled are not known in advance. On a page cache miss, np.take faults their 16 rows’ pages one at a time at QD1, putting 2-4.6 ms on the critical path. This was why p90 was 4-5 ms above p50.

Right before the gather, one process_madvise call now requests all 16 rows’ pages (the first and last page of each row). The faults overlap and the reads go out in parallel. The generic function from the prefetch thread (unique, then merge ranges) took 0.15 ms, so this is a dedicated version that reuses its iovec array. p90 at 8k went from 16 to 12.4-13.0 ms, the mean from 72-74 to 80-82 tok/s.

3.5 Memory: reserved grew by 190 MiB with allocated unchanged

The first version used +190 MiB on GPU1 when idle and reached 24,126 MiB after a prefill (the ceiling is 24,124). torch.cuda.memory_allocated and its peak were exactly the same as the baseline; only reserved was higher.

The cause was a change in the allocation pattern during load_model. Because the hyper-connection weights were no longer repacked by Marlin (no more allocate-new, free-old), the holes left by other layers’ repacks moved, and unusable fragments stayed behind. torch.cuda.empty_cache() does not return them (fragments smaller than expandable segments’ 2 MiB granularity).

The fix restores the original allocation pattern. The hyper-connection weights and scales are cloned once (the same allocate-new, free-old as Marlin). For INT6, allocating the two planes separately left 100-190 MiB of fragments, so both go into one allocation of the original size and then the original is freed. The GPU peak with 4 x 247k resident went from [24072, 24100, 24090] to [24022, 24090, 24038] MiB, lower than before.

3.6 Results

Same harness, same compose settings as production:

The contribution of each, switched by environment variables in one build:

On short prompts, 03 and 04 each give only about 0.5 ms alone, and 2.5 ms together. Host-side processing and GPU time were alternately the bottleneck, so improving only one is limited by the other.

The remaining 10 ms is almost all GPU time: about 3.2 ms of graph per rank, plus lm_head and sampling.

For correctness: the hyper-connection fusion matches an fp32 reference exactly at one token; the MoE path, on real weights over 192 real decode steps, differs from the original by at most 9.7e-3 relative (a few BF16 ULPs) with no change in expert selection; the INT6 planes pass a bit-exact round-trip test; and server-side logprobs differ by 0.001995 old vs new against 0.002010 old vs old, about the same as run-to-run variation.

3.7 Drive-side latency

Decode sometimes stopped for 100-780 ms. During a stop, the threads’ wchan was folio_wait_bit_common (waiting for page I/O). Measuring 4 KiB O_DIRECT random reads continuously, the NVMe holding the table (a DRAM-less QLC drive) has periods of several seconds where every read takes about 140 ms (970 ms at worst). Another NVMe (a Crucial P310) under the same conditions had a p50 of 0.106 ms and a max of 22 ms.

Software cannot fix that, so I eventually moved the table to the P310. The drive for the PLE table should be chosen by the tail (p99 and max) of its random-read latency.

4. Approaches tested and not adopted

Three approaches were implemented and measured, then not adopted.

4.1 TP=3

Last time I chose PP over TP because the dimensions do not divide by 3. With padding, TP=3 is not impossible. I measured the performance upper bound first, to decide whether it was worth implementing.

Weights were --load-format dummy, and the config was rewritten into divisible shapes (attention heads 24 -> 36, KV heads 2 -> 3, linear attention heads 16 -> 18, shared expert 640 -> 768, vocab to a multiple of 192), the same amount of work as after padding. Routed experts were split by EP into 171 / 171 / 170.

The PP=3 baseline was measured under the same dummy conditions. Both used only four of the release patches (vllm.patch, decode-01, decode-02, ttft-01), without decode-03 and later from section 3. QSA offload only works with TP=1, so it was off and the KV was in VRAM (max-model-len 32k, prefix caching off, no PLE). The PP=3 numbers below are therefore not directly comparable to the production numbers in section 3.

The kernels that read weights shrink to a third, but the roughly 2,000 small kernels per token still all run on each of the three cards. On top of that come about 96 all-reduces per step. vLLM’s custom all-reduce does not support a world size of 3 (Supported world sizes: [2, 4, 6, 8, 16]), so it falls back to NCCL. Prefill is limited by communication on three cards connected only by PCIe.

The upper bound is +10%, and the QSA host offload would also need changes for TP, so it was not implemented.

4.2 MTP experts in host memory

This model comes with a 4B MTP module (a draft head for speculative decoding). Last time I left it out for lack of VRAM. What if MTP’s routed experts were quantized to INT4, kept in pinned host memory and read by the GPU directly over UVA?

Short prompts are faster; every other metric is worse.

An expert row on the host reads top-10 x 2.46 MB = 24.6 MB over PCIe, 0.9 ms per row. Verification is 2 rows per request, so +1.8 ms per decode step. In prefill, the draft head runs its MTP layer on every prompt row (to build the KV of its own QSA layer), so each 512-row chunk reads all 512 experts over PCIe, +45 ms. Four streams lose because with parallel requests the three PP stages are already busy, and verification doubles the rows while the UVA reads land on rank 2 alone.

Weights can be kept in host memory only when they are used for the computation of many tokens. In decode, expert weights are used for 1-4 tokens, so the PCIe transfer time is not amortized.

4.3 A RAM cache for the PLE table

Would keeping the most-used rows of the table in RAM cut cold reads? I built an 11.8M-token trace from the official benchmark generations, Japanese documents and the vLLM source, turned it into rows with the server’s hash, and compared LRU, a static hot set and the page cache in a simulator written in C.

The trace touched 41.3M rows = 12.3 GiB, or 71.6 GiB as pages. A row cache is 12.8 times denser than the page cache. But a third of the trigram rows are seen for the first time, and no cache holds those. With the top 8M rows (2.4 GiB) as a static set, cold NVMe reads drop by 32% in prefill and 18% in decode. The share of decode steps that read at least one row from NVMe goes only from 0.46 to 0.40. Since decode-06 issues the 16 reads in parallel, a step with one read waits about as long as a step with sixteen.

In time, that is about 10% of cold TTFT and under 1% of decode. Together with the problem of how to build and distribute the hot list, it was not adopted.

5. Decode, third round: caching K/V rows in idle VRAM

In 3.1, the extra 2 ms at long context was the PCIe read of QSA’s K/V, which does not overlap with compute. The approach here is to reduce the reads themselves.

Instrumenting the indexer showed that consecutive steps select 60-85% of the same rows (0.72 overlap with the previous step on average, 0.86 with the union of the last 4 steps). Keeping the rows read earlier in VRAM would avoid most reads.

The problem is VRAM allocation. Even one request needs about 20 MiB per rank, and in the worst case of 4 x 262k, GPU1 has 24 MiB free.

The VRAM comes from the prefill staging arena

In section 5.3 of the last article, prefill copies the context into a VRAM work buffer (the staging arena, 4 QSA blocks = 4 x 12,144 x 2 KiB ≈ 95 MiB per GPU) one section at a time before computing on it. The arena is used only inside a prefill forward, and it is idle during decode. It is allocated at startup, so reusing it costs no extra VRAM.

• The arena is split by the number of QSA layers on the rank (4) into one 12,144-row direct-mapped cache per layer. One slot holds one token’s K and V for both KV heads, 2 KiB.
• The key is the physical cache slot (block x block_size + offset), not the logical position. Requests that share blocks through prefix caching share rows too.
• The host K/V only changes on a store, so dropping a slot’s row right after the store keeps the cache consistent. Reassigned blocks also go through a store.
• A staged prefill overwrites the arena, so all tags are dropped after it.

One layer, one step

1. invalidate: drop the rows of the slots just written, and add 1 to the rank’s epoch.
2. plan: a column that hits writes hitmark = epoch into its slot. A column that misses claims its slot with claim = atomic_max(epoch << 32 | slot).
3. attention: a kernel based on the split-K kernel that reads hits from VRAM and misses from host memory. A miss writes its row and tag only if its claim won and no column hit that slot in this step.

Races are avoided by two rules. A slot that was hit is never written in the same step (with hitmark == epoch nobody gets the right to write it). The tag of a slot that may be written is never read (with hitmark != epoch it counts as a miss). So reading and writing tags inside the same kernel is safe, and the only kernel needed for grid synchronization is the plan. claim and hitmark carry the epoch, so they need no clearing between steps.

Pinned host memory and VRAM share one address space under UVA, so one load handles both hits and misses:

ptr = tl.where(hit, region_ptr, host_ptr)
row = tl.load(ptr)

One load reads VRAM on a hit and host memory on a miss. With separate loads for hits and misses the kernel needed 108 KB of shared memory, over the 3090’s 101 KB limit, and did not launch.

A hit reads the same bytes as the host copy and the tile arithmetic is unchanged, so the output is bit-identical, not approximate. 1,000 steps of synthetic data (a deliberately small region with many hash collisions, host rows rewritten between steps, the arena overwritten, and so on), and a check on the real model in eager mode that runs both paths and compares with torch.equal more than 8,000 times per rank: zero mismatches.

Why it stops at 12 rows

The first version had no row limit, and GPU1 ran out of memory during a 4-stream 8k prefill. It compiled and loaded a Triton kernel variant for a large eager batch in the middle of serving, and that memory was not budgeted. GPU1 has 34 MiB free.

Up to 12 rows the kernel has a single launch configuration, which is loaded when the cudagraphs are captured at startup. Only decode-shaped batches (12 rows or fewer) use the cache; larger batches use the old path (and their stores still invalidate).

Results

The hit rate is 86-91% on real text and 88-97% on synthetic prompts, higher than the plain overlap with the previous step (0.72) because rows from a few steps back are still in the region. For the kernel alone: 174 µs per layer reading everything from the host, 78 µs at 72% overlap, 38 µs at 100%. At 0 overlap it is 177 µs, almost no overhead.

Long-context decode is now about as fast as short-prompt decode.

6. Adding image input

Until now it ran text-only (--language-model-only). The model includes a ViT (27 blocks, width 1152, 0.836 GiB in BF16).

It does not fit as is / unused weights

Just dropping --language-model-only makes vLLM build visual on every PP rank: 0.84 GiB x 3. Rank 0 already peaks at 24,072 MiB, and it does not fit even with the whole KV budget removed.

vLLM’s runner runs the encoder on the first rank only, so the ViT on the other ranks is unused. While checking this, I found that embed_tokens (INT8, 0.60 GiB) and lm_head (0.60 GiB) were on all three ranks too. Only the first rank uses embed and only the last rank uses lm_head. The upstream qwen4_exp code does this, and does not use PPMissingLayer the way other vLLM models do.

mem-01 fixes only that. The weights loaded per rank went from [21.58, 21.55, 21.55] to [20.98, 20.44, 21.05] GiB (0.60 / 1.11 / 0.50 GiB less on ranks 0 / 1 / 2). 2.2 GiB across the three cards had been used by unused weights.

Streaming the ViT from host memory

Even without the duplicates, keeping the ViT on rank 0 had two problems.

1. Out of memory during loading. The MoE Marlin repack temporarily takes 400 MiB. If the ViT is already on the GPU, rank 0 fails there. -> The ViT is built and loaded on the CPU and moved to the GPU in process_weights_after_loading(), after all repacks are done.
2. Resident, it forces the KV budget from 783 down to 480 MB. 262,144 x 4.00 becomes 2.44.

So the 27 ViT blocks stay in pinned host memory and are streamed block by block, with prefetch, into two slots on the GPU (58 MiB). While block i computes, block i+1 is copied into the other slot on a side stream. Each block’s forward is wrapped so that, right before it runs, its parameters’ .data is swapped for a view into the slot (the ViT blocks are @support_torch_compile modules whose __call__ goes straight to forward, so nn.Module hooks never fire).

As described in 4.2, weights can be kept in host memory when they are used for many tokens. ViT weights are used for all patches of an image (thousands), so this condition holds.

Copying one block over PCIe 4.0 x16 takes 1.16 ms regardless of the number of patches. Compute scales with patches, about 4 ms per block at 1 MP. When compute is longer, only the first block’s copy is visible. The break-even is about 0.4 MP; smaller images are copy-bound, but the difference is at most about 12 ms. TTFT is mostly the language model’s prefill, so for a 2 MP image it stayed at 0.85 -> 0.85 s.

Images are scaled down to 2 MP (2,048 tokens) with --mm-processor-kwargs '{"max_pixels":2097152}'. Without a limit, the startup profiling assumes a 16.7 MP (16,384-token) image and eats rank 0’s headroom.

With images on, every rank peaks lower than it did text-only. With 4 x 247,823 resident: needle 4/4, and 45/45 images alongside new long prefills.

7. Putting MTP on the GPU

The approach in 4.2 was not adopted, but mem-01 freed VRAM, so I reconsidered it.

With 0.5-1.1 GiB freed per rank, the PP layer split can change from an even 16 / 16 / 16 to 16 / 17 / 15. On 09-15, a rank with 17 layers ran out of memory because it carried the unused embed and lm_head (mem-01 reduces this rank 1’s weights by 1.11 GiB). Now rank 2 has about 2.1 GiB free at its peak. With MTP’s routed experts in INT4 g128 (1.17 GiB) plus its dense part (0.17 GiB), it should fit on the GPU.

What MTP under PP needed (mtp-01)

In stock vLLM, enabling MTP with PP>1 did not even start.

• Build the draft head’s embed_tokens with the quantization config. Under PP>1 the target’s embed is on the first rank and cannot be shared, so the draft head keeps its own. It was built without quant_config, so the INT8 embed could not be loaded into it at all.
• Fix a branch in forward. The draft head exists only on the last rank and receives the target’s hidden states directly. But it branched on whether the global PP rank is the first one, so under PP>1 it always took the intermediate-tensors path, which nobody fills, and an assert fired.
• Count the draft head’s layers in the QSA host RAM estimate. The draft head’s attention is also full attention and allocates a host pool. The estimate counted only the target’s layer_types.
• Size the attention block for the widest rank. The draft head’s QSA layer makes the last rank’s QSA group 5 layers instead of 4. The block size is shared by all ranks, so if it is not sized for this group, the group’s page exceeds the mamba page and every block in the pool grows.

Fitting it on three cards (mtp-02)

Two more tensors that do not need to be on the last rank were moved.

• The draft head’s lm_head: the draft head built a quantized copy of lm_head, only for it to be replaced right after loading by the target’s lm_head, which is on the same rank. For a copy that only lives until it is replaced, the Marlin repack took 0.6 GiB during the load, and rank 2 failed. It is now a PPMissingLayer from the start.
• The draft head’s embed_tokens (INT8, 0.60 GiB): a step gathers only a few rows from it, so it is built and loaded on the CPU and left in pinned memory behind a UVA view after loading. A PCIe transfer of a few rows is negligible.

make_mtp_int4.py quantizes the MTP experts to INT4: RTN, g128, symmetric, with a per-group search for the clip that minimizes squared error. It does not modify the downloaded checkpoint; it writes three files to a separate directory, which compose mounts on top.

Results

This table used the development configuration (--kv-cache-memory-bytes 783000000, --max-num-seqs 4), measured on the same day with the same script, switching only MTP on and off. The 3x3090+MTP column in the first table is a separate run with the release settings.

With the experts on the GPU, the problems from 4.2 (1.5x prefill, UVA reads) are mostly resolved. A single stream is 15-43% faster. The slowdown at four streams comes from speculative decoding itself: when GPU utilization is high, doubling the verified rows makes steps slower.

The capacity drops not because of VRAM but because R from section 2 (the number of blocks one 262,144-token request takes) goes from 61 to 85.

• QSA: the draft head’s QSA layer makes rank 2’s QSA group 5 layers, which shrinks the block size from 12,144 to 9,776 tokens.
• Linear attention: with speculative decoding, vLLM adds num_speculative_tokens (1) page per mamba group for the draft token. With the 2 pages of align mode, that is 3 pages per group.
• Block size: the gated delta-net conv state grows by one draft token (20,480 B), so the mamba page becomes 3,227,648 B. The block size 9,776 is the largest multiple of 16 for which 5 layers x 66 B/token (the VRAM-resident part of QSA) fits in that page.

At 783 MB that is 242 blocks / 85 = 2.85 requests; at 550 MB, 170 blocks / 85 = 2.00.

The decode-07 row cache is split differently too. The staging arena holds “QSA layers on the widest rank x block size” tokens, which in the MTP configuration is 5 x 9,776 x 2 KiB ≈ 95 MiB. Rank 2 uses it as 5 layers x 9,776 rows; ranks 0 and 1 split the same-size arena into 4 layers of 12,220 rows.

The release uses --kv-cache-memory-bytes 550000000 (262,144 x 2, 41.2 GiB pinned) and --max-num-seqs 2. When it accepted four requests, an image request came in as the third while two 250k requests were resident, triggering 20 preemptions and a 177-second wait for one image.

MTP is a trade-off, so both configurations are distributed. For one user who wants latency, 3x3090+MTP; for throughput and context, 3x3090. My production runs the latter.

8. Structured output returns 500 with MTP + PP (mtp-03)

With the MTP build and qwen-code, replies were returned, but 500 errors occurred in the background.

ERROR [backend_xgrammar.py:168] Failed to advance FSM for request chatcmpl-... for tokens 198.
ERROR [scheduler.py:2073] Unexpected: grammar rejected tokens [198, 198] for request chatcmpl-.... Terminating request.
"POST /v1/chat/completions HTTP/1.1" 500 Internal Server Error

What failed was not the main reply but a side request that qwen-code sends alongside it with response_format: json_schema (input suggestions and the like).

It happens only with MTP and two or more running requests. My patches touch neither the scheduler nor the grammar bitmask, so this is an upstream vLLM problem.

The cause is DraftTokensHandler in the V2 model runner. With structured output requests present, the engine defers sampling, receives the GPU’s drafts through take_draft_token_ids(), and builds the grammar bitmask from them. But DraftTokensHandler returns only the drafts of the last sampled batch.

Under PP, the V2 runner lets a request decode only every pp_size steps. With two requests, they alternate between batches. So the drafts of the request whose bitmask is being built do not come back and stay -1. When the scheduler sees -1, it masks the bonus position (the token after the draft) with the grammar state from before the draft. The GPU, meanwhile, verifies and accepts the real draft. At the bonus position, a token the grammar forbids after the accepted draft can then get through, and accept_tokens rejects it and ends the request. It matches that the rejected tokens come in pairs of [draft, bonus].

The fix makes DraftTokensHandler keep the latest drafts of each request and return all of them. These are the drafts the GPU verifies next, so the scheduler and the GPU agree. Finished requests are removed from the dict. Before the fix, 9 of 16 requests at 2 parallel returned 500; after it, all 45 were OK, with the same acceptance rate.

The conditions are “V2 runner + async scheduling + PP>1 + speculative decoding + structured output + 2 or more parallel requests”.

9. Quality: reproducing the official benchmarks

Last time I only reported KLD (0.0149 nats against BF16). KLD measures how far the distribution moved, not whether the answers are right. From the table on the official model card, I ran the three benchmarks that need neither an agent environment nor an external judge, with the same image and weights as production. The server was the production compose with only prefix caching turned off (--max-num-seqs 4, 262,144 x 4).

All three are within about 1 SE of the official numbers; a single run shows no degradation from quantization. Sampling used the generation_config.json defaults (T=1.0 / top_p=0.95 / top_k=20) with thinking on and no output limit. The longest response was 201k tokens on LiveCodeBench; a limit of 81,920 truncates GPQA and LiveCodeBench answers. Generation totaled 9.7M tokens over about 13 hours (about 205 tok/s on average).

Measuring decode speed on the same server (400 output tokens) gave 95-100 tok/s with one stream and 283-298 tok/s total with four. This is not directly comparable to the 243-245 tok/s in the first table: prefix caching differs (with it on, the mamba align mode kernels from section 3.1 are added) and so does the measurement script.

Eight streams were slower, at 230 tok/s total. For that run 8 was added to cudagraph_capture_sizes, so 8-token steps also ran as cudagraphs. The decode-04 / 05 fused kernels, however, handle up to 4 tokens, so 8-token steps take the unfused path. I did not switch the fused kernels on and off at 8 streams, so their share of the difference is not isolated. The benchmarks ran with four streams (cudagraph_capture_sizes [1,2,4]).

LiveCodeBench’s grader has problems. As of 2026-09, testing_util.py in lcb_runner fails 8 correct solutions.

Qwen3.8 models often write sys.stdin.buffer.readline and sys.stdout.buffer.write, so plain lcb_runner gives 86.3. Regrading the stdin problems in a real subprocess gives 92.4. Every problem that passed in lcb_runner also passed in the subprocess, so the difference is not from a more lenient check. LiveCodeBench scores across models are comparable only with the same grader.

10. The current configuration

--pipeline-parallel-size 3 --tensor-parallel-size 1
--max-model-len 262144 --max-num-seqs 4
--max-num-batched-tokens 512
--enable-prefix-caching --mamba-cache-mode align
--gpu-memory-utilization 0.97 --kv-cache-memory-bytes 783000000
--limit-mm-per-prompt '{"image":4,"video":0}' --mm-processor-kwargs '{"max_pixels":2097152}'
--compilation-config '{"mode":0,"cudagraph_mode":"FULL_DECODE_ONLY","cudagraph_capture_sizes":[1,2,4]}'

VLLM_PLE_MMAP_PATH=<table file>   VLLM_QSA_KV_OFFLOAD=1
VLLM_MEMORY_PROFILER_ESTIMATE_CUDAGRAPHS=0
PYTORCH_CUDA_ALLOC_CONF=expandable_segments:True
cap_add: SYS_PTRACE   ulimits: memlock -1

Compared with last time, VLLM_QSA_KV_OFFLOAD_MAX_GIB and VLLM_QSA_KVO_ARENA are gone. The first because the host pool allocation now checks free RAM itself, the second because the shipped value is now the default. All patches are applied by the Dockerfile at build time, so the release and my production run byte-identical code.

Summary

• TTFT: PLE table reads were synchronous at queue depth 1. Prefetching the whole prompt with process_madvise at admission brought cold TTFT from 44 to 8.3 s.
• KV capacity: prefix caching changes the blocks per request from 42 to 61. A block’s cost is mostly pinned host RAM; with exact-size pinned allocation, 262,144 x 4 fits.
• Decode: the host-side serial processing under PP, the low effective bandwidth of small GEMVs, PLE page cache misses and the PCIe read of QSA K/V were addressed in turn, taking 1 stream from 59-73 to 95-99 tok/s and 4 streams from 120-154 to 243-245 tok/s.
• Memory: the prefill staging arena is reused as a K/V row cache during decode, and the ViT and MTP are placed in the space freed by removing unused embed / lm_head. All three cards are within a few tens of MiB of their limit, so no new VRAM is allocated.
• Not adopted: TP=3 (+10% upper bound, −60% prefill), MTP experts in host memory (−34% at 4 streams), a RAM cache for the PLE table (about −10% TTFT, under −1% decode).

The patches and compose files are all here. The target is vllm/vllm-openai 0.29.1rc1.dev47+gdc36fcce9. https://huggingface.co/Minachist/Qwen3.8-Flash-Next-INT4-Mixed-AutoRound