Eval bug: qwen4exp (Qwen3.8-Flash-Next) CUDA: decode slows linearly with context

[lukolszewski]

Name and Version

$ /app/llama cli --version
version: 0.4.0-dev (build 10909, commit 44b47fd77)
built with GNU 12.3.0 for Linux x86_64

Operating systems

Linux

GGML backends

CUDA

Hardware

Ryzen 7950x on x670 MB 192MB RAM 5x RTX3090 (3 of them via USB4, but this detail is not relevant to the bug)

Models

unsloth/Qwen3.8-Flash-Next-GGUF UD-Q4_K_XL -ngl 99, layer split across 5 GPUs,
-ot per_layer_token_embd.weight=CPU -c 524288 --parallel 2

Problem description & steps to reproduce

This is a likely a duplicate of a previous report erroneusly closed by reporter #27856, I can't reopen it so I'm submitting this one with all my detail. The issue #27856 is not a duplicate of #27856, that one appears to be fixed (it reports much more severe deterioration).

When runnig the PR #27742 branch I initaly noticed substantial performance degradation about halving the decode speed starting at 54 tok/s every 50k context size or so.

Using the latest master df03399 this is much improved to:

• 37.2tok/s to 50.8tok/s at 42k ctx (+37%)
• 22.5tok/s 48.0tok/s at 111k (+113%),
• 11tok/s to 41.8 tok/s at 250k

but still deteriorating to "nighly speed" on my hardware at the bottom of the context.

Please note both generation is affected (severely) as well aas prefill (a lot less). Prefill almost halves the speed at max 260k context so far less impactful.

I let my LLM run over the code to see what improvements it finds and by the morning indeed it managed to make changes that helped a lot to:

• 37.2tok/s to 50.8tok/s at 42k ctx (+37%)
• 22.5tok/s to 48.0tok/s at 111k (+113%)
• 11tok/s to 41.8tok/s at 250k (+280%)

It performed some correctness validation that suggests we didn't brake anything, but I'm far from claiming a "fix" at this stage (have yet to validate properly).

The diff is almost 2k lines over 7 patches so it is too big to include here. I will provide it upon request.

First Bad Commit

it was built from pull/27742/head

Relevant log output

I also include a brief LLM written sumamry of changes made below:

qwen4exp decode depth-decay — LLM generated "fix" summary

Arch: qwen4exp (Qwen3.8-Flash-Next, 125B-A6B; 12 QSA sparse-attn + 36 Gated-DeltaNet layers, indexer.top_k=2048, compress_ratio=4).
Problem: single-stream decode slows ~linearly with context depth even though QSA is supposed to bound attention to a fixed top-k budget. Root cause (nsys-profiled): the top-k path was running as a dense attention scan over all n_kv, plus per-token re-pooling of the whole cache and MBs of host→device + P2P input traffic.
Base: llama.cpp master df03399b8, branch levers. Build/run: CUDA 12.4.1, sm_86, image llama.cpp:lever2b-cuda124; KV_TYPE=f16 is mandatory (quantized KV can't reach the sparse MMA kernel at batch 1 on Ampere).

Results (single-stream decode, t/s, 5×RTX 3090)

37.2tok/s to 50.8tok/s at 42k ctx (+37%)
22.5tok/s to 48.0tok/s at 111k (+113%)
11tok/s to 41.8tok/s at 250k (+280%)

(shallow essentially unchanged: 54.4 → 56.5. Prefill improves as a side effect: 609 → 716 t/s @111k.)

Changes (7 patches, llama-levers/000*.patch)

 lever1     src/models/qwen4exp.cpp                 |  9 +++++----   (+ incidental Dockerfile pip fix)
 lever3     src/llama-memory-hybrid-idx.{cpp,h}     |  99 ++++++      block-granular top-k
            src/models/qwen4exp.cpp                 |  72 ++++
 sparse256  ggml/src/ggml-cuda/fattn-mma-f16.cuh    |  3 +           whitelist DKQ=DV=256 MMA
 lever2     src/llama-memory-hybrid-idx.{cpp,h}     | 590 ++++++     persistent pooled block-key cache
            src/llama-kv-cache.{cpp,h}, qwen4exp.cpp
 lever2b    src/llama-memory-hybrid-idx.{cpp,h}     |  59 ++++       device-resident block-member cells
            src/llama-kv-cache.{cpp,h}, qwen4exp.cpp

lever1 — enable the sparse-FA path (the big one). The enablement PR #27742 shipped top-k as a dense mask with a // TODO: enable sparse attention. Un-TODO it: pass top_k->ne[0] as n_kv_max so flash-attn treats the finite mask entries as a sparse K/V set — O(top_k) instead of O(n_kv). Drops attention from ~11 ms to 0.22 ms/token @250K.

-    // TODO: enable sparse attention when we are ready
-    ggml_tensor * cur = build_attn_mha(q, k, v, nullptr, kq_mask_top_k, nullptr, nullptr, 0, kq_scale, il);
+    // LEVER1: passing top_k->ne[0] as n_kv_max -> sparse K/V set (deepseek4/GLM-DSA pattern, #27970)
+    ggml_tensor * cur = build_attn_mha(q, k, v, nullptr, kq_mask_top_k, nullptr, nullptr, top_k->ne[0], kq_scale, il);

sparse256 — instantiate the kernel for this head shape. may_use_sparse only whitelisted the MLA shapes; qwen4exp is DKQ=DV=256, ncols2=8 and that MMA template instance already exists. One clause:

-    return (DKQ == 512 && DV == 512 && ncols1 == 1 && ncols2 == 8) ||
+    return (DKQ == 256 && DV == 256 && ncols1 == 1 && ncols2 == 8) ||
+           (DKQ == 512 && DV == 512 && ncols1 == 1 && ncols2 == 8) ||
            (DKQ == 576 && DV == 512 && ncols1 == 1 && ncols2 == 16);

lever3 — block-granular top-k. Scores are block-constant and the budget is whole blocks, so select blocks first and expand only the ~513 selected blocks to cell indices (via blk_cells), instead of expanding block scores to all n_kv cells and sorting those. Removes the O(n_kv) expansion get_rows + one O(n_kv) mask add and shrinks the sort by compress_ratio. Fully-future blocks get −inf bias; the incomplete tail block's cells ride a small per-token side input.

lever2 — persistent pooled block-key cache. The indexer re-gathered, mean-pooled, RMS-normed and re-RoPEd all cached blocks every token though only the newest changed. Add a third per-QSA-layer buffer of pooled+normed+roped block keys; each ubatch recomputes only blocks it dirtied (host-side structural compare of member cells — also catches defrag/rollback/seq-ops with no invalidation hooks) and ggml_set_rows them in. Dirty count is pow2-bucketed (≥4) for CUDA-graph topology stability.

lever2b — device-resident block-member cells. After 1–3 the residual depth cost was input traffic (~3.1 MB/token HtoD + ~13.2 MB/token P2P for blk_cells/mask). Store each block's member cells in the block cache's V tensor (f32; ggml_set_rows takes no i32) and gather the top-k expansion on-device — removes the blk_cells term entirely.

Validity

Every patch was A/B'd against the immediately-prior build and kept only if it passed all gates, all at temperature 0:

• Exact needle-in-haystack retrieval at 63k and 111k (a random code string embedded mid-context must be reproduced verbatim) — guards against the sparse path silently dropping the right K/V.
• Multi-turn prefix reuse + rewind and 2 concurrent slots — guards against the persistent block-key cache going stale across defrag/rollback/seq ops.
• Throughput read from timings.predicted_per_second on POST /completion, cache_prompt:false, n_predict:48, ignore_eos:true.

A comparable dense model on the same rig (arch=qwen35, plain GQA, no indexer) stays flat at these depths, confirming the decay was QSA/indexer-specific, not KV bandwidth.

Status / upstream

Upstream #28012 (this exact CUDA decay) is closed. Patches are self-contained and PR-able in order: lever1+sparse256 (smallest), then lever3, then lever2+lever2b (shared machinery). Remaining depth losses: host-side O(n_kv) grouping scan in set_input_qsa; per-GPU f16 mask fan-out; flash_attn_mask_to_sparse_indices re-scanning the mask per layer (an indices-in FA API would remove both). Full detail + nsys attribution in new-issue.md / issue-speed.md §9.

[lukolszewski]

I've been using the code with the LLM written "fix" and I noticed it only worked if I used 16 bit KV cache. With q8 kv the problem persisted so I've let the LLM work on that problem too. The "fix" is on my branch (linked above this comment). I write "fix" in quotes, because I haven't reviewed the code. I can't vouch for its correctness, but I hope the diff will be helpful to resolve this comprehensively.

[svgop]

Following this with great interest — we're serving Qwen3.8-Flash-Next on a single RTX 5090 at 256K context and your numbers match the pathology we've independently profiled (QSA attention cost dominating at depth). Could you share the patch set you offered (lever1+sparse256, lever3, lever2+2b)? Happy to test the full series on our rig (5090/CUDA 13.x, Blackwell sm_120) and report measured decode/prefill at 42K/111K/250K with and without. The self-contained PR-able landing order you described sounds ideal.

[wussh]

Independent reproduction on a single A100 MIG slice (IQ3_XXS, native 256k, f16 KV) + two checks that isolate what the decay is not

Different hardware, different quant, same pathology. Rig: one A100 80GB PCIe MIG 7g.80gb slice (sm_80), llama.cpp 0.4.1-dev build 11011 (aa39d7a3e, 2026-09-17), ISTA-DASLab/Qwen3.8-Flash-Next-GSQ-RCO-GGUF IQ3_XXS + BF16 mmproj, -ngl 99 -c 261888 -np 1 -ub 1024 --jinja -fa auto -lm mmap -lzm on, single stream.

Fresh prime + 256-token decode burst on the same prefix:

Context	Prefill tok/s	Decode tok/s
32k	1 196	46.5
64k	1 054	39.1
128k	846	29.3
192k	719	23.3
224k	668	21.1
250k	638	20.1

Linear fit: t_decode ≈ 17.3 ms + 0.13 µs × ctx ⇒ at 250k, ~32.5 ms of the ~50 ms step (≈65 %) grows with context while the context-free part stays ~17 ms. The slice is saturated and the host is idle during both phases (SM 1.34–1.40 GHz, HBM 1 512 MHz max, 269–282 W, CPU ~1.0 core of 16, pod RSS 9–12 GB) — consistent with your nsys attribution (dense O(n_kv) top-k/attention scan), not KV byte traffic.

1) Quantising KV makes it worse (your KV_TYPE=f16 requirement, reproduced from the other direction)

Config	Decode @128k	Decode @250K	Prefill @128k/@250K
KV f16 (default)	29.3	20.1	846 / 638
-ctk q8_0 -ctv q8_0	22.3 (−24 %)	14.1 (−30 %)	850 / 641
-ctk q8_0 (V f16)	19.7 (−33 %)	12.7 (−37 %)	851 / 642

The quantisation was active (VRAM −2 956 MiB) and tool calls stayed valid (p50 2.0 s), so this is the kernel path leaving the FA/MMA kernels — the same constraint you describe as "quantized KV can't reach the sparse MMA kernel at batch 1 on Ampere". A fix that only works with f16 KV is exactly what these numbers predict.

2) CUDA graphs are not the decay — A/B with GGML_CUDA_DISABLE_GRAPHS=1:

Metric	Graphs on	Graphs off
decode @128k (256 tok)	29.3	26.2 (−10 %)
decode @250K (256 tok)	20.1	18.6 (−8 %)
prefill @128k / @250K	846 / 638	860 / 644
first-forward init	100.6 s	101.1 s

Replay is active and worth keeping (~8–10 % on decode), but the context-proportional loss is not capture/launch overhead: the one-time ~100 s first-forward init is unchanged with replay disabled (it is CUDA module/kernel setup, not capture).

3) Both gates are still closed in this pin (verified against the source of build 11011, not just from the issue text):

// src/models/qwen4exp.cpp:764-767 — the sparse variant sits right above, commented out
//ggml_tensor * cur = build_attn_mha(q, k, v, nullptr, kq_mask_top_k, nullptr, nullptr, top_k->ne[0], kq_scale, il);
ggml_tensor * cur = build_attn_mha(q, k, v, nullptr, kq_mask_top_k, nullptr, nullptr, 0, kq_scale, il);

// ggml/src/ggml-cuda/fattn-mma-f16.cuh:1760
static constexpr __host__ __device__ bool ggml_cuda_flash_attn_ext_mma_f16_may_use_sparse(
        const int DKQ, const int DV, const int ncols1, const int ncols2) {
    return (DKQ == 512 && DV == 512 && ncols1 == 1 && ncols2 == 8) ||
           (DKQ == 576 && DV == 512 && ncols1 == 1 && ncols2 == 16);
}

and build_qsa_top_k (qwen4exp.cpp:542-687) expands block scores to all n_kv cells before ggml_top_k. This model is DKQ = DV = 256, ncols1 = 1, ncols2 = 8 (n_embd_head_k/v = 256, n_embd_k/v_gqa = 512), so the sparse MMA kernel is not even reachable here. lever1 + sparse256 are, concretely, a one-line call-site change plus a one-line whitelist clause.

We can canary the series on this rig (A100 MIG sm_80, IQ3_XXS, f16 KV, 261 888 ctx, single stream) and report decode/prefill at 32k/128k/192k/250k, needle retrieval at 10/50/90 % depth, tool calling and VRAM — an independent, differently-shaped validation on the arch family where f16 KV matters. If nobody has landed the first two patches yet, we are happy to prototype them ourselves and open a PR against upstream.

[Atomic-Germ]

Doesn't literally every model do this?

[lukolszewski]

Following this with great interest — we're serving Qwen3.8-Flash-Next on a single RTX 5090 at 256K context and your numbers match the pathology we've independently profiled (QSA attention cost dominating at depth). Could you share the patch set you offered (lever1+sparse256, lever3, lever2+2b)? Happy to test the full series on our rig (5090/CUDA 13.x, Blackwell sm_120) and report measured decode/prefill at 42K/111K/250K with and without. The self-contained PR-able landing order you described sounds ideal.

Yes the patches are on my fork branch issue/28734 https://github.com/lukolszewski/llama.cpp/tree/issue/28734 commit b22bc5b

However I discovered a day later the fix is working only fas long as kv is f16 and I wanted to use q8 to save VRAM so I have next level of fixes on my branch q8-compact https://github.com/lukolszewski/llama.cpp/tree/q8-compact commit 4d6ef5a

I was going to recommend you use the q8-compact if you use quantized kv, but today I discovered another two bugs that are not covered by this issue and require a longer writeup which I'll post tomorrow(as it is late now). In short, when we serve this model with multiple slots, for example slot 0,1,2 and user in the middle (slot 1) is idle, but slots 0 and 2 are generating there is an over 50% slowdown because (as I understand it now) the graph is not reused at all, but regenerated for each token. I have a "fix" for that too on branch server-seq-compact https://github.com/lukolszewski/llama.cpp/tree/server-seq-compact latest commit now (b75e380).

I suggest to use that one to test. Also I built my images for Blackwell arch, as one of my cards is 5060 ti 16GB so I can confirm it works fine on Blackwell.

The other problem I do not have fixes for is(I'm not sure if this one is fixable):

• when prefill is running on one slot and one or more slots generate they will get an atrocious rate of 1-3tok/s :-( why? Because the entire ubatch takes some time to run on prefill (1s for ubatch of 1024 on my hardware) and only inbetween there is an opportunity to generate. This can be made a little bit better by lowering ubatch, but then you kill your prefill processing speed.... Considering I run a lot of RAM caching and slot saving, prefill takes couple seconds tops so I can live with that for now.

Doesn't literally every model do this?

A little bit maybe, but not to the extent seen here, I compared to Qwen3.8-27B (that uses different architecture) and the slowdown was tiny, on the order of a couple t/s from 1k to 260k context.

So, this is a new architecture, it is expected to have issues. I'm glad we have it in llama.cpp at all so quickly. On my hardware vllm is a nogo (I have to use layer paralell and I need hand placed layers and other items) so llama.cpp is the only game in town.

[goodbadwolf]

With speculative decoding on, the merged sparse path helps prefill and costs decode, and one clause separates them

Rig: one RTX PRO 6000 Blackwell (sm_120), unsloth/Qwen3.8-Flash-Next-GGUF UD-Q4_K_XL snapshot 38bb39ee, -ngl 999 -ot per_layer_token_embd=CPU -fa on -np 1 -c 524288 --rope-scaling yarn --rope-scale 2 --yarn-orig-ctx 262144 -ctk q8_0 -ctv q8_0, default -b 2048 -ub 512, MTP speculative decoding at --spec-draft-n-max 4. Baseline is 53b1389d0b; candidate is master 0c3626ec06 (b11070, which carries #28770 and not #28536) with PR#28243's MTP commits merged on top, clean merge. I reached lever1 and sparse256 independently before finding this thread, so first: both reproduce exactly as @lukolszewski and @wussh describe. Everything below is about speculation, which nobody here appears to be running.

Read decode as steps per second, from the log rather than from acceptance. With MTP, tok/s is the step rate times the tokens a step returns, and the second factor is draft acceptance, which is not controlled. The server already prints what is needed: mean len on its draft acceptance line is tokens per step, so a request's step rate is its own tok/s divided by its own mean len. One example from the runs below, three repetitions of one prompt at one depth: 41.80, 53.44 and 42.86 tok/s, a 28 percent spread, and 16.85, 16.91 and 16.74 steps per second, a 1.0 percent spread. Across eleven server starts the within-configuration step-rate spread stayed between 0.2 and 3.1 percent. 1 + n_max x acceptance agrees with mean len to under 1.5 percent and works as a fallback, but n_max is the maximum draft length rather than the length drafted, so the logged value is the one to trust.

The merged sparse path splits by phase. Prefill tok/s and decode steps/s, three depths, fresh server per configuration, three repetitions, same prompts in the same order, decode taken per request from mean len.

Depth	baseline 53b1389d0b	master, sparse disabled	master as merged	master, with the Q->ne[1] guard
9,411	1962.3 / 45.20	1978.4 / 45.82	1982.9 / 46.17	1977.6 / 45.90
106,572	1586.9 / 32.35	1590.0 / 31.63	1683.8 / 30.88	1677.6 / 31.78
350,734	989.7 / 17.01	994.3 / 16.97	1201.4 / 16.40	1194.0 / 16.85

Column 2 is the control: with the sparse path disabled inside shall_use_sparse, the candidate binary lands within 2.3 percent of the old baseline on decode at every depth and within 0.5 percent on prefill. So the effect in column 3 is the sparse path and not the 34 intervening commits: +5.9 and +20.8 percent of prefill, against -2.4 and -3.4 percent of decode. The decode cost grows with depth, which the phase argument below predicts.

Why the two phases disagree. shall_use_sparse computes n_gather = (ncols1 == 1 ? Q->ne[1] : ncols1) * n_kv_max and engages at K->ne[1] >= max(4096, 2*n_gather), and switch_ncols1 tries the ncols1 == 1 variant first. n_kv_max here is 2051, from min(n_kv, indexer_top_k + r - 1) with top_k 2048 and compress_ratio 4. A prefill flash-attention node sees one physical ubatch, 512 by default, so its n_gather is 1,050,112 and that branch cannot fire at any depth this model reaches. Prefill falls through to ncols1 = 8, ncols2 = 8 and the union kernel, whose own sparse check needs 2*8*2051 = 32,816. An MTP verify batch has Q->ne[1] = 5, so n_gather is 10,255 and the threshold is 20,510: decode takes the one-column path from 20k of context onward. The 9,411 row sits below that and does not move.

That row is a weaker control than it looks. It shows the 34 commits cost nothing at a depth where the sparse path is inactive on both axes, which is not the same as showing they cost nothing at depth. The QSA_SPARSE_MODE=0 column is the control that carries the argument.

The cost looks like the index build rather than the attention. I filed #29281 separately with test-backend-ops numbers. On the 11 sparse configurations it generates, all of them at nb=1, #28770 costs 64 percent median and 113 percent worst. Dividing the added time by the mask width gives a constant, 0.589 to 0.594 ns per column across a twelve-fold range of kv and across n_kv_max 512 and 2048, which points at flash_attn_mask_to_sparse_indices rather than at the attention kernel. Its new inner loop has a runtime trip count and a data dependent exit, so the #pragma unroll around it no longer yields straight line code, and at ncols1 == 1 the body always runs once.

If that is right, the scan fix belongs before the guard. A per-column cost is paid whatever the batch width, so a 5-token verify batch spreads it over five queries and a 1-token decode over one, which would make 3 percent here and 64 percent there the same mechanism. Specialising ncols1 == 1 might remove enough of it that the guard is unnecessary. I cannot say that it would, because the f16 conversion below is a second cost with the same shape and I have not measured it. I would fix the scan and re-measure the guard against that, rather than take both.

Forcing the verify batch onto the union kernel instead is much worse, so the current dispatch choice is right. Skipping the ncols1 == 1 branch so a 5-token batch reaches ncols1 = 8 gives 45.70 / 24.88 / 11.59 steps/s at the three depths, against 46.17 / 30.88 / 16.40 as merged. Prefill was unchanged at 1984.9 / 1686.4 / 1199.5, confirming only the decode path moved. The loss grows with depth.

What does work is refusing the sparse path for narrow batches. Column 4 adds one clause to shall_use_sparse, if (Q->ne[1] < 64) return false;, which covers both call sites. That clause has been serving this unit since 2026-09-21, behind an environment lever so the three states can be compared from one binary. Against the dense control it keeps +5.5 and +20.1 percent of prefill and returns decode to within +0.5 and -0.7 percent, inside the repeat spread at every depth.

For continuations at depth the benefit depends on the append size, and small appends lose. Production prefill on this unit is overwhelmingly short continuations rather than cold prompts, so this is the case that decides whether the guard is worth having. Two append sizes at the same 349k depth, same build pair, deployed against guarded: 893-token appends take 2053 and 1765 ms, so +16.3 percent, while 119-token appends take 885 and 912 ms, so -2.9 percent. Reading those four points as a fixed cost plus a per-token cost, the guard cuts the per-token cost 27 percent, from 1.51 to 1.10 ms, and adds about 75 ms of fixed cost, with break-even near 184 tokens. Two points per line is a weak fit and I would not defend the coefficients, but the sign change is measured rather than modelled. My best statistic for this unit's real traffic is a median of 287 newly evaluated tokens per call, which the model puts at about +4 percent. So the honest summary is a large win on cold deep prefill, a few percent on a median continuation, and a small loss on very short ones.

I would guess the added fixed cost is the top-k index build over the whole cache, which runs whatever the append size, but I have not measured that.

The q8_0 restriction may be deeper than kernel reachability. The batch-1 story is that quantized KV selects the VEC kernel, which has no sparse path. Speculation removes that: Q->ne[1] = 5 is above the VEC condition, so MMA_F16 is selected and the sparse path is reachable at q8_0. It is reachable and it still loses for decode, because MMA_F16 sets need_f16_K/need_f16_V and fattn-common.cuh converts the whole cache with to_fp16(K_data, K_f16, ggml_nelements(K), ...) before any index is read. That conversion is O(n_kv) and unconditional, so it caps what the sparse path can win, and a 5-token batch is too narrow to amortise it. I want to be careful about how far that goes: the dense MMA_F16 path pays the same conversion, so this limits what the sparse path can win without demonstrating why it loses, and moving the conversion after the gather would not on its own prove the guard unnecessary. But if the reading is right, q8-compact and a Q->ne[1] guard are pushing on the same constraint from two ends.

Separately, PR#28213's gather cannot fire alongside MTP at all. Its condition requires n_tokens to equal n_stream, and a five-token verify batch against one stream never matches. Measured: decode steps/s of 45.30 / 32.00 / 17.05 with it enabled against 45.61 / 31.62 / 17.22 disabled, identical within the repeat spread at all three depths. With #28770 merged it looks redundant anyway.

What I have not shown. Retrieval of three planted constants at 10/50/90 percent depth scored 6/6 at 106,630 and 6/6 at 350,792 as merged, and the guarded configuration scored 6/6 at 350,792 against the running server. Six trials against one planted set is a smoke test, not a qualification of either path. I also had a weighted "percent of total serving time" figure here and have removed it: the journal window it rested on does not reproduce, so the honest claim is the per-phase numbers above and the continuation measurement, not a single blended number.

Happy to re-run any of this against a scan fix, to test server-seq-compact, or to open the guard as a PR if it still earns its place afterwards.

[albertnsoliz]

Hey, found an improvement for this, unsure if this is helpful but discussion/further validation would help:
#29326

[lukolszewski]

There have been some rewrites in the master of the files involved so I was curious if the latest master is any better overall. Also is it worth upgrading and porting my patches to get some improvement. Sadly on my hardware the answer is it is not (unless someone is running single user maybe )

I had to put the MTP head on a slow GPU (rtx5060 ti 16GB) because 5 slots of 262k context fills my 6x rtx3090s to the brim and the potential slowdown is really only relevant in single user scenario where if I dropped users to 3-4 I'm guessing I could maybe get to 60tok/s for a single user with a tiny prompt.... but there is still the speed collapse with longer prompts and the MTP helps very little in a multi slot setup.

So for now I'm staying on my patched image for this model.

Here are my results after porting the remaining patches with and without mtp (master without any changes was worse at the moment):

Setup. Qwen3.8-Flash-Next UD-Q4_K_XL (unsloth), 6× RTX 3090, layer split -ts 0.85,1,1,1,1,1, 5 slots × 262k ctx, q8_0 KV, -fa on, -ub 512 -b 2048, llama-server. "old" = master df03399b8 (2026-09-10) + the patches from this thread (sparse-FA levers, compact-gather decode, server seq compaction, prefill admission, pipeline parallelism). "new" = master bed0a8566 (2026-10-02) + the same patches ported on top of its k-pool indexer (the levers that upstream had meanwhile re-done in its own way were dropped; the compact-gather decode was re-done on the k-pool selection). MTP = unsloth's self-contained mtp-Qwen3.8-Flash-Next-Q8_0.gguf as --spec-type draft-mtp --spec-draft-n-max 2, the head on an otherwise idle RTX 5060 Ti (--spec-draft-device), temperature 1.0. Decode numbers are tokens/s per stream at the given context depth, measured on the same prompts with the same scripts, back to back, on the same day. Prefill is a single 82k-token request.

	4k	58k	117k	221k	3 users @58k, each	5 users @58k, each	prefill 82k
old, ub512 (production)	51.7	47.0	45.6	44.3	25.1	18.7	1260
new, ub512	48.5	39.5	35.8	31.2	18.5	12.4	348
new, ub256 (pipelining fits)	53.9	48.3	46.6	44.3	23.5	15.6	1226 on an empty box, 628 once the other slots hold context
new + MTP (head on the Ti)	47.7	40.3	46.4*	27.4	19.8	13.5	356

* 92 % draft acceptance on this repetitive benchmark prompt; everywhere else acceptance was 0.4–0.6 at temperature 1.0.

• At -ub 512 the new tree cannot reserve its pipeline-parallel compute buffers (the k-pool indexer's scoring tensors are several GB at 5 × 262k), falls back to no pipelining, and in that state decodes 7–30 % slower and prefills at a quarter of the speed. At -ub 256 pipelining fits and the new tree matches the old one for a single user at every depth.
• Multi-user it stays 6–17 % behind, and its prefill halves once the other slots hold context: the k-pool indexer scores every stream's pools against every token of the batch (the foreign ones are masked afterwards), so the per-token cost grows with the total context held across all slots rather than with the request's own.
• Stock master aborts at start-up with this MTP head (the MTP block is dense, compress_ratios[48] == 0, but the graph still builds the k-pool inputs and set_input hits an unallocated tensor); that needed a fix before anything could be measured. Two more fixes were needed to make MTP usable at all at depth: the k-pool layout was rebuilt from scratch (O(n_kv) on the host) after every tail edit, and speculative decoding edits the tail on every step (221k: 15.9 t/s before, 27.4 after).
• Even then MTP never beats the old tree without MTP, except at the one point where the benchmark prompt is drafted almost perfectly. At 3–5 concurrent users it loses 20–30 %, since the draft passes cost the same ~2 × 4 GB of weight reads per step regardless of how many tokens they carry. With the head on a 3090 instead of the Ti I would expect roughly +10–15 % for a single user at shallow context and break-even around 60k, but not a change for the depth or the concurrency picture. Speculative verification (3 tokens per stream per step) also runs on the dense O(n_kv) attention path, which is where the remaining depth penalty with MTP comes from.

So for this workload (several users, 50k–250k contexts, quantized KV) the September tree plus the patches above remains faster, and MTP is not worth the upgrade on this hardware.

[lukolszewski]

For anyone who wants to test: the patches now live in a separate repo with binaries and container images,
https://github.com/lukolszewski/llama.cpp-multigpu (release multigpu-20261006: tarballs for CUDA 12.9 covering
V100 → RTX 5090, and CUDA 13.4 for Ampere+, plus ghcr.io/lukolszewski/llama.cpp-multigpu images; CUDA only, as
that is my optimisation target).

Measured on 6 × RTX 3090 (PCIe, no NVLink), Qwen3.8-Flash-Next UD-Q4_K_XL, same command line on both sides
(5 slots × 262144, q8_0 KV, -fa on, -b 2048 -ub 512, layer split). Baseline = master df03399b8, the commit
the patches branch from — my earlier wording "compared with the original PR branch" was wrong.

workload	context	master df03399b8	llama.cpp-multigpu	gain
prefill, 1 session	50k	568 t/s	2127 t/s	3.7×
prefill, 1 session	250k	263 t/s	2111 t/s	8.0×
prefill, 5 sessions (aggregate)	250k	265 t/s	2060 t/s	7.8×
generation, 1 session	50k	25.7 t/s	41.9 t/s	1.6×
generation, 1 session	250k	10.2 t/s	33.7 t/s	3.3×
generation, 5 sessions (per session)	250k	2.3 t/s	27.3 t/s	10.8×

Upstream's rate falls with context and collapses under concurrency; the patched build stays roughly flat
(~2100 t/s prefill at every depth, 27–46 t/s per session of decode). Full 20-row grid (5k/50k/150k/200k/250k ×
1 and 5 sessions), protocol, raw JSON and server logs: repo README and benches/.

Two things to keep apart, because this issue is about one of them:

1. The decode slowdown with context depth reported here — the first two patches (block-granular indexer top-k
   with a persistent pooled-key cache, and a compact-gather attention path so a quantized KV cache keeps the sparse
   speedup). That is the "generation, 1 session" rows above.
2. Multi-GPU layer split on PCIe without NVLink — several sequences decoding without the GPUs waiting on each
   other, pipeline-parallel prefill, prefilling one slot while others decode, async prompt-cache transfers. That is
   where the prefill and 5-session numbers come from, and it is a different problem from this issue; I have not
   filed it separately.

Caveats: one machine, one model; single-GPU or NVLink systems may not benefit and some patches could regress them.
Mixed prefill + decode is better than before but still the weak spot. The code was written with an LLM and
validated by measurement and output checks (needle-in-haystack, temp-0 output identical between paths), not by
review; I am not going to open PRs from it myself. Each change is one identifiable commit described in
docs/multigpu/patches.md, and the series is attached to the release as git format-patch output, so anyone is
welcome to pick up any piece.

Separately, n-gram lookup speculation (prompt lookup decoding, no draft model) roughly doubles single-user decode on
code rewrites here (44 → 111 t/s); it costs other users a graph re-plan per draft round, so the fork caps it at two
concurrently drafting users by default.

[am17an]

This should be fixed now