spec : add DFlash2 support (local convolution + candidate selector)

[SubSir]

Overview

DFlash2 adds two more modules compared to DFlash, grouped dynamic depthwise convolution and candidate selector.

The formula is for reference. out[i,c] = Σ_t (base[t,c] + δ[i,t,g(c)]) · x[i−t,c]. base is the static kernel, δ is the dynamic part predicted from the input.x is divided into groups g(c) and δ is shared across a group. Another one is edge(p→c) = ⟨A[p] ⊙ project(h), B[c]⟩ + unary[c]. A/B are the codebooks for the predecessor and the candidate. This position's hidden is projected by project(h)to the same rank to multiply A. And unary[c] is the draft's own score for c.

DFlash2 is enabled when the checkpoint is DFlash2; no need to use extra flag.

We evaled Qwen3.8-27B at Q4_K_Mon Apple M5 Pro 64 GB. And the test data is first 8 problems of GSM8K.
Concurrency is set to 1 and temperature is 1.0, top-p 0.95, top-k 20. We use default xhigh reasoning and set
2,048 maximum new tokens.

Method	Draft quant	Block	Decode TPS	Decode speedup	Acceptance	Normal stops
Autoregressive	—	—	10.42	1.00×	—	8/8
DFlash2	BF16	8	19.31	1.85×	4.92	8/8
DFlash2	Q8_0	8	18.43	1.77×	5.08	8/8
DFlash2	Q4_K_M	8	18.89	1.81×	5.03	8/8

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: YES. Using AI to inspect and analyze code in the llama.cpp project, and using AI to update the code.

[ggml-gh-bot]

Hi @SubSir, thanks for your contribution!

Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:

• AI-generated content: While code is allowed to be generated by AI, please write the PR description and commit messages on your own without the help of AI.

---

Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below.

[SubSir]

I've revised the PR description. It's fully hand-written now.

[showgood163]

Hi, there.

After some testing in with qwen 3.8 in llama.cpp webui, I found out that vision support (with PDFs) is broken, while the native mtp is not.

checkpoint from https://huggingface.co/incoai/Qwen3.8-27B-DFlash2-GGUF
parameters (as suggested):
model-draft = Qwen3.8-27B-DFlash2-Q4_K_M.gguf
spec-type = draft-dflash
spec-draft-n-max = 7

[KaruroChori]

Very impressive for single parallelism: on my b70 I get almost a 3x performance boost under my testing scenarios with 5 tokens (regression for higher values).
But the moment I increase the number of agents, it quickly drops to 1 tok/s or something like that. MTP has no such behaviour.
Is this reproducible in other architectures or is it something specific to intel?

[Shamish]

Vision requests fail with draft-dflash: M-RoPE positions rejected by the draft context

Tried this PR on a Tesla V100 (sm_70, CUDA 12.9) with Qwen3.8-27B-UD-Q4_K_XL as target,
incoai/Qwen3.8-27B-DFlash2-GGUF:Q4_K_M as drafter and the matching mmproj-F16.
Text generation and tool calling work fine. Any request carrying an image dies:

srv  send_error: decode() failed: failed to process speculative batch
process: llama_decode(ctx_dft) failed rc=-1 (n_tokens=4, offset=0)

This is not an OOM — the model loads at 31015 / 32768 MiB, text works on the very same
instance, and the peak while processing the image is 31699 MiB.

Cause

For a multimodal prompt the target batch is M-RoPE: batch_in.pos holds position tuples
(n_pos_per_embd == 4) and an image span repeats the temporal component across its tokens.
The DFlash drafter is a plain qwen3 model with n_pos_per_embd == 1, and for such a
context llama_batch_allocr requires strictly continuous positions per sequence
(seq_pos_min(s) == p0 + 1, src/llama-batch.cpp). process() copies the target
positions verbatim:

batch_inject.pos[i] = batch_in.pos[j];

so the injection batch inherits repeated/jumping positions and llama_decode(ctx_dft)
rejects it with rc=-1.

Note that common_speculative_impl_draft_mtp::process() never hits this because it bails
out on embedding batches early — the // TODO: how to make it work with vision tokens?
right above that check is, I believe, the same underlying issue.

Suggested fix

The draft model keeps its own KV cache and only needs the row order to match the target;
the absolute position values do not have to. Numbering the injected rows densely per
sequence fixes vision and is a no-op for text-only batches, where the target positions are
already dense:

             batch_inject.n_tokens = n_chunk;
             std::memcpy(batch_inject.embd, inp_g, (size_t) n_chunk * n_embd_dec * sizeof(float));
+            // A multimodal target batch is M-RoPE: batch_in.pos holds position tuples and an
+            // image span repeats the temporal component. ctx_dft has n_pos_per_embd == 1 and
+            // llama_batch_allocr then demands strictly continuous positions per sequence, so
+            // copying batch_in.pos verbatim makes llama_decode(ctx_dft) reject the batch with
+            // rc=-1 and drops vision requests. The draft keeps its own KV cache and only needs
+            // the row order to match, so number the injected rows densely per sequence. For a
+            // text-only batch the target positions are already dense and this is a no-op.
+            std::vector<llama_pos> pos_next(n_seq, -1);
             for (int32_t i = 0; i < n_chunk; ++i) {
                 const int32_t j = offset + i;
                 GGML_ASSERT(batch_in.n_seq_id[j] == 1);
                 const llama_seq_id seq_id = batch_in.seq_id[j][0];
                 GGML_ASSERT(seq_id >= 0 && seq_id < (llama_seq_id) n_seq);
-                batch_inject.pos[i]       = batch_in.pos[j];
+                if (pos_next[seq_id] < 0) {
+                    pos_next[seq_id] = llama_memory_seq_pos_max(llama_get_memory(ctx_dft), seq_id) + 1;
+                }
+                batch_inject.pos[i]       = pos_next[seq_id]++;
                 batch_inject.n_seq_id[i]  = 1;
                 batch_inject.seq_id[i][0] = seq_id;
                 batch_inject.logits[i]    = false;

Verification

Built 5ecbe1ac1 with and without the patch (-DGGML_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=70),
same server flags, temperature 0, seed 42:

Case	Before	After
Image prompt	rc=-1, request dropped	described correctly
Text	63 t/s, acceptance 103/132	63 t/s, acceptance 103/132
Text right after an image prompt	n/a	62.5 t/s, acceptance 132/184

The third row is the one I was most worried about — the draft KV survives a multimodal
prefill, acceptance stays in the normal range, and no rc=-1 or CUDA errors appear in the log.

Caveat: only tested single-sequence (parallel = 1) on one target/drafter pair, so a
second opinion on the multi-sequence path would be welcome.

Unrelated data point: --spec-draft-n-max 7 is bad on Volta

The README of the GGUF repo suggests --spec-draft-n-max 7, which matches the H200
numbers. On a V100 it is a sizeable regression; a sweep over three prompts (3 prompts x 2
runs, mean TG) gave:

n-max	2	3	4	5	7	built-in MTP (n4)
TG, t/s	51.9	53.1	55.5	52.2	44.9	52.1

At n-max 7 a long generation drafted 1356 tokens for 315 accepted and dropped to 27 t/s
against 39.8 for the built-in MTP. Q8_0 of the drafter was consistently slower than
Q4_K_M while costing ~0.9 GiB more VRAM, matching the acceptance-length table in the
model card.

---

Written by Claude (Anthropic's Claude Code), posted from the account of the human who ran
the hardware. The investigation, the patch and every number above come from actual runs on
the V100 box described at the top — nothing here is inferred or estimated. Please review
the patch on its merits; I can rerun any check you would like to see.

[treo]

@Shamish There is already a fix for the reported problem at z-lab#1

As for the Performance on a single RTX3090, it seems to be not significantly better than MTP:

MTP:

$ ./build/bin/llama-server -m ~/.cache/huggingface/hub/models--unsloth--Qwen3.8-27B-GGUF/snapshots/f1bfb127c64f7072bdd2cad55f258b9c8b2910fe/Qwen3.8-27B-UD-Q4_K_XL.gguf  --fit off -np 1  -fa on --spec-type draft-mtp --spec-draft-n-max 3 --reasoning-preserve -c 32000 -ctk q8_0 -ctv q8_0

$ python speed_bench.py --url localhost:8080 --model qwen3.8-27b --concurrency 1 --limit 2
Warning: You are sending unauthenticated requests to the HF Hub. Please set a HF_TOKEN to enable higher rate limits and faster downloads.
speed_bench: loaded 22 samples from bench=qualitative category=all
speed_bench: 100%|██████████████████████████████████████████████████████████████████████████████████████████| 22/22 [11:57<00:00, 32.60s/sample]

Summary (elapsed=717.20s)
category       samples  avg_prompt_t/s  avg_pred_t/s  avg_latency  accept_rate
-------------  -------  --------------  ------------  -----------  -----------
coding         2        454.24          59.84         69.160s      0.6047     
humanities     2        508.66          57.24         56.283s      0.6075     
math           2        27.86           55.76         17.593s      0.5927     
qa             2        85.46           54.37         18.066s      0.5589     
rag            2        704.60          59.45         40.403s      0.6777     
reasoning      2        60.51           54.47         18.138s      0.5954     
stem           2        28.90           54.39         17.867s      0.5981     
writing        2        612.05          61.74         42.117s      0.6609     
multilingual   2        164.13          65.08         4.339s       0.7778     
summarization  2        131.48          54.36         4.965s       0.5473     
roleplay       2        397.24          58.95         69.665s      0.6814     
overall        22       288.65          57.79         32.600s      0.6317  

DFlash2 (using maxn=4 because 7 was slower):

$ ./build/bin/llama-server -m ~/.cache/huggingface/hub/models--unsloth--Qwen3.8-27B-GGUF/snapshots/f1bfb127c64f7072bdd2cad55f258b9c8b2910fe/Qwen3.8-27B-UD-Q4_K_XL.gguf  --fit off -np 1  -fa on --spec-type draft-dflash --spec-draft-n-max 4 --reasoning-preserve -c 32000 -ctk q8_0 -ctv q8_0 -md ~/.cache/huggingface/hub/models--incoai--Qwen3.8-27B-DFlash2-GGUF/snapshots/6cb5872e2cee6b4e780a8414922350be8e42d65c/Qwen3.8-27B-DFlash2-Q4_K_M.gguf

$ python speed_bench.py --url localhost:8080 --model qwen3.8-27b --concurrency 1 --limit 2
Warning: You are sending unauthenticated requests to the HF Hub. Please set a HF_TOKEN to enable higher rate limits and faster downloads.
speed_bench: loaded 22 samples from bench=qualitative category=all
speed_bench: 100%|██████████████████████████████████████████████████████████████████████████████████████████| 22/22 [11:40<00:00, 31.85s/sample]

Summary (elapsed=700.65s)
category       samples  avg_prompt_t/s  avg_pred_t/s  avg_latency  accept_rate
-------------  -------  --------------  ------------  -----------  -----------
coding         2        271.99          68.00         48.928s      0.6691     
humanities     2        470.70          59.87         62.873s      0.5610     
math           2        28.45           60.14         29.131s      0.5827     
qa             2        82.64           53.74         13.088s      0.5039     
rag            2        701.27          66.42         37.589s      0.6378     
reasoning      2        58.79           60.55         29.046s      0.5827     
stem           2        29.38           60.20         28.899s      0.5827     
writing        2        607.63          69.16         38.621s      0.6552     
multilingual   2        167.16          62.26         5.469s       0.5316     
summarization  2        136.16          59.60         4.324s       0.5690     
roleplay       2        375.62          64.43         52.359s      0.6758     
overall        22       266.35          62.21         31.848s      0.6153  

[Stoney49th]

A test with 2 / 4 / etc. concurrenc (-np) would be nice - MTP has some issues there currently and if this also scales with parallel requests, that would be a totally different story :)

[ruixiang63]

Thanks for the PR. It would be great if you could run the benchmark on SpeedBench using the corresponding checkpoints and report the performance numbers. @SubSir

[KaruroChori]

A test with 2 / 4 / etc. concurrenc (-np) would be nice - MTP has some issues there currently and if this also scales with parallel requests, that would be a totally different story :)

I tried them on an intel B70, but it was a no go. Massive drop to basically baseline performance with 2, and drop to the 1 tk/s range with 4. By comparsion MTP has a less pronounced degradation.

[Zeldatales]

Tested this branch on AMD Strix Halo (gfx1151) with the Vulkan backend. Builds clean and the DFlash2 drafter loads and runs correctly. Sharing numbers since this looks like an untested configuration — it's an integrated GPU on unified memory, which behaves quite differently from discrete VRAM.

Setup

SoC	AMD Ryzen AI MAX+ 395 (Strix Halo), Radeon 8060S, gfx1151
Memory	128 GiB unified LPDDR5X — GTT 62.5 GiB, dedicated VRAM 0.5 GiB
OS	Arch Linux, kernel 7.1.8, mesa 26.1.7 (RADV)
Build	this PR @ 5ecbe1ac1 (build 10498), -DGGML_VULKAN=ON, -march=native
Target	ggml-org/Qwen3.8-27B-GGUF Q4_K_M (17.7 GB)

llama-server -m Qwen3.8-27B-Q4_K_M.gguf -md <drafter>.gguf \
             --spec-draft-n-max 4 --no-mmap -ngl 999 -ngld 999 --ctx-size 8192

Measured over /v1/chat/completions, max_tokens=300, temperature=0, 5 samples per cell, end-to-end wall clock (includes HTTP + prompt eval, so absolute numbers sit below server-side generation rates — relative comparison is unaffected).

Results (tok/s ± σ)

Drafter	Size	Prose prompt	Code prompt
DFlash2 Q8_0	1.9 GB	18.96 ± 0.62	22.68 ± 0.69
MTP Q4_0	1.6 GB	18.36 ± 0.88	20.15 ± 0.54
DFlash2 BF16	3.6 GB	17.77 ± 0.78	20.31 ± 0.52
no drafter	—	11.24	—

Observations

1. Gain over MTP is workload-dependent. DFlash2 Q8_0 is +12.6% over MTP on code generation (clears 3–4σ) but only +3% on prose (inside noise). A single prose prompt would have shown the two as equivalent — worth noting for anyone benchmarking this.

2. The smaller drafter wins here. BF16 trails Q8_0 by ~9–10% on both prompts. On unified memory the drafter and target contend for one memory channel, and the extra 1.7 GB of weight traffic appears to cost more than the improved acceptance rate returns. I'd expect this to invert on a discrete GPU.

3. block_size saturates at n_extract, and is not reachable from the CLI. Sweeping it on the Q8_0 drafter (code prompt, 7 samples each):

block_size	tok/s	σ
3	19.27	0.58
4	21.29	1.02
5	21.91	1.02
8	21.91	1.02

Throughput climbs to block_size 5 and then stops. block_size 8 matched block_size 5 sample-for-sample across seven deterministic runs (21.47/21.70/21.11/21.12/23.68/21.31/23.01 vs 21.39/21.71/21.12/21.13/23.69/21.32/23.02). The startup line explains it — n_extract=5 never moves, matching the drafter's dflash.block_count=5. So anything above the drafter's depth is a no-op, while anything below it throttles drafting and costs real throughput.

Two things follow. It might be worth clamping or warning when block_size > n_extract, since the setting silently does nothing. And there is currently no way to set it without editing the file: it lives only in GGUF metadata, and --override-kv applies to the target model — there's no --spec-draft-override-kv. I patched the drafter GGUF in place to run this sweep.

4. Minor: a [spec] failed to measure draft model memory: failed to create llama_context from model warning appears during startup, but loading proceeds and inference is correct.

For context on the platform

The ROCm/HIP backend on the same machine (packaged b10485, MTP drafter) runs 29% slower than Vulkan — 13.37 ± 0.01 vs 18.78 ± 0.42 tok/s on prose. So these Vulkan numbers are the fast path on this hardware, not a fallback.

Limitations

Two prompts, five samples per cell (seven for the block_size sweep). Acceptance rate is content-dependent and the prose/code gap is already 12 points, so other workloads may land elsewhere. The block_size sweep used the code prompt on the Q8_0 drafter only.

Happy to run additional configurations if useful.

[jhstatewide]

but not sure if this is all just nonesense

Typically, this type of thing doesn't help. Now you're presenting a thing that you don't even know works and someone has to waste time trying to figure out if it even works? You're just adding noise and no signal.

It's great that you want to help out with OSS projects. Just find a different way to contribute.

[gabrielcosi]

Blackwell numbers (sm_120).

RTX PRO 4000 Blackwell SFF, 24 GB, 70 W cap, CUDA 12.8. Built at 5ecbe1a with -DCMAKE_CUDA_ARCHITECTURES=120.

Target Qwen3.8-27B NVFP4, draft incoai/Qwen3.8-27B-DFlash2-GGUF Q4_K_M, --ctx-size 32768 --kv-unified -fa on -ctk q8_0 -ctv q8_0 --spec-draft-n-max 7. Each row fires N concurrent requests capped at 500 tokens. Aggregate is total tokens over wall time.

aggregate t/s	np=1	np=2	np=4
no speculation	22.3	38.9	52.7
draft-mtp, n-max 2	38.8	56.7	106.2
draft-dflash, n-max 7	38.6	95.7	140.0

Separately, a single 1100-token code generation at --ctx-size 131072: 57.4 t/s for DFlash2 against 44.4 t/s for MTP.

[Stoney49th]

Blackwell numbers (sm_120).

RTX PRO 4000 Blackwell SFF, 24 GB, 70 W cap, CUDA 12.8. Built at 5ecbe1a with -DCMAKE_CUDA_ARCHITECTURES=120.

Target Qwen3.8-27B NVFP4, draft incoai/Qwen3.8-27B-DFlash2-GGUF Q4_K_M, --ctx-size 32768 --kv-unified -fa on -ctk q8_0 -ctv q8_0 --spec-draft-n-max 7. Each row fires N concurrent requests capped at 500 tokens. Aggregate is total tokens over wall time.
aggregate t/s np=1 np=2 np=4
no speculation 22.3 38.9 52.7
draft-mtp, n-max 2 38.8 56.7 106.2
draft-dflash, n-max 7 38.6 95.7 140.0

Separately, a single 1100-token code generation at --ctx-size 131072: 57.4 t/s for DFlash2 against 44.4 t/s for MTP.

Thats kind of strange. In my setup with either running it with/without layer split on R9700's (2x) i dont see the "near linear" scaling for multiple sessions. np=1 is higher with ~45-48TPS, but it dops to 23/24 per slot for np=2, and falls below no speculation for long contexts and long generations (above 130k). Thats on vulkan though - I need to retest ROCM. Will try the PR the upcoming days on the dual card config...

[darksylinc]

I am (non-exhaustively and non-scientifically) testing on R9700 + Vulkan.

At spec-draft-n-max = 4 it produced the best results, I see a small bump over MTP (around 49-60 t/s, but averages at around 52 t/s) with some peaks at 74 t/s (fastest I've ever achieved with this model). MTP produces around 40-60t/s but will often sit at around 46-48 t/s.

Those gains are not as spectacular as claimed, but are an improvement. My settings (MTP settings also included, commented out):

./llama-server --jinja --port 8085 -fa on -np 1 --log-colors on --no-webui --metrics --models-max 1 --no-warmup --models-preset models/config.ini

model = AI/LLM/Models/Qwen3.8-27B-UD-Q4_K_XL.gguf
ctx-size = 131072
ctk = iq4_nl
ctv = iq4_nl
ub = 2048
b = 1024
temp = 1.0
top-p = 0.95
top-k = 20
repeat-penalty = 1.0
presence-penalty = 0.0
dev = Vulkan0
chat-template-file=AI/LLM/Models/qwen3.8_chat_template.jinja
#spec-default =
#spec-type = draft-mtp
#spec-draft-n-max = 4
fit = off
spec-type = draft-dflash
spec-draft-n-max = 4
model-draft = AI/LLM/Models/Qwen3.8-27B-DFlash2-Q4_K_M.gguf

The JINJA is froggeric's v22.2

[Stoney49th]

I am (non-exhaustively and non-scientifically) testing on R9700 + Vulkan.

At spec-draft-n-max = 4 it produced the best results, I see a small bump over MTP (around 49-60 t/s, but averages at around 52 t/s) with some peaks at 74 t/s (fastest I've ever achieved with this model). MTP produces around 40-60t/s but will often sit at around 46-48 t/s.

Those gains are not as spectacular as claimed, but are an improvement. My settings (MTP settings also included, commented out):

model = AI/LLM/Models/Qwen3.8-27B-UD-Q4_K_XL.gguf
ctx-size = 131072
ctk = iq4_nl
ctv = iq4_nl
ub = 2048
b = 1024
temp = 1.0
top-p = 0.95
top-k = 20
repeat-penalty = 1.0
presence-penalty = 0.0
dev = Vulkan0
chat-template-file=AI/LLM/Models/qwen3.8_chat_template.jinja
#spec-default =
#spec-type = draft-mtp
#spec-draft-n-max = 4
fit = off
spec-type = draft-dflash
spec-draft-n-max = 4
model-draft = AI/LLM/Models/Qwen3.8-27B-DFlash2-Q4_K_M.gguf

The JINJA is froggeric's v22.2

Could you maybe also try with two parallel requests in np=2 or np=3? would be nice to have a second reference if the scaling is an issue on my side...

[darksylinc]

Could you maybe also try with two parallel requests in np=2 or np=3? would be nice to have a second reference if the scaling is an issue on my side...

I'm not familiar with client-side apps that make parallel requests 🫤
I could try 2/3 manually crafted curl commands.

But if you've got suggestions on how to try parallel requests I'll gladly take it.

[Stoney49th]

Could you maybe also try with two parallel requests in np=2 or np=3? would be nice to have a second reference if the scaling is an issue on my side...

I'm not familiar with client-side apps that make parallel requests 🫤 I could try 2/3 manually crafted curl commands.

But if you've got suggestions on how to try parallel requests I'll gladly take it.

https://github.com/ggml-org/llama.cpp/tree/master/tools/server/bench/speed-bench

think thats the easiest and most straight forward. Also tried to reproduce gabrielcosis results on ROCM and did a second MTP run for sanity, but nowhere near. The dropoff at higher fills above 100k is even heftier than vulkan, down to ~18TPS per slot for two filled to 110k each. Vulkan stays at around 22TPS then, scales to 3x15TPS or the same fill degree. Still need to establish a reliable baseline before the Dflash2 test makes sense...

[ThConMan]

Thanks for this PR. I built it to try DFlash2 on Windows and hit a load failure I could not resolve, so here is a precise repro in case it is useful.

Summary: with this branch, the z-lab DFlash2 drafter fails to load when paired with a target, with invalid vector subscript. Notably this is progress over master, which fails earlier with a tensor-count mismatch, so the v2 layout is clearly being recognised.

Environment

build 10498 (5ecbe1ac1), this PR, clean checkout
Windows 11, MSVC 19.44.35227.0, CUDA 13.3
cmake -B build -DCMAKE_BUILD_TYPE=Release -DGGML_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=120
RTX 5070 Ti Laptop (sm_120, 12 GB)

Repro

llama-server.exe -m Qwen3.8-27B-Q2-XYZ-v2.gguf \
  --spec-type draft-dflash \
  --spec-draft-model Qwen3.8-27B-DFlash2-Q4_K_M.gguf \
  -c 8192 --parallel 1 -ctk q8_0 -ctv q8_0 -fa on

common_speculative_init_result: loading draft model 'Qwen3.8-27B-DFlash2-Q4_K_M.gguf'
llama_model_load: error loading model: invalid vector subscript
common_speculative_init_result: failed to load draft model
srv    load_model: failed to load draft model

Drafter: z-lab/Qwen3.8-27B-DFlash2-GGUF.

What I ruled out

• Not the drafter file. Loading it standalone gets through model load and fails only at context creation with dflash requires ctx_other to be set, which is expected for a drafter.
• Not the quantisation. Q4_K_M and Q8_0 fail identically.
• Not an unusual target quant. I first hit this with a third-party 2-bit quant, then reproduced it identically with bartowski/Qwen3.8-27B-GGUF IQ2_XXS.
• Not out-of-range tap points. The drafter declares dflash.target_layers = [6, 20, 34, 48, 62], all valid against the target's 65 blocks.
• Not tensors indexed by target layer id. Every per-layer create_tensor in src/models/dflash.cpp uses the drafter's own loop index.

I also chased common_speculative_init_result calling llama_model_load_from_file(params.model.path...) while logging params.speculative.draft.mparams.path, which looks like a mix-up — but common_base_params_to_speculative sets result.model = params_spec.mparams first (common/speculative.cpp:2382), so it is correct. Mentioning it only so nobody else loses time on the same false lead.

One possibly useful detail

MSVC's std::vector::at() throws std::out_of_range with exactly the text invalid vector subscript. So this is an .at() overrun rather than an operator[]. On GCC/Clang the same site would either report differently or, if it is an operator[], read out of bounds without complaint — which may be why this has not shown up in Linux testing.

The drafter's own metadata, in case a shape assumption is involved:

dflash.block_count = 5          dflash.embedding_length = 5120
dflash.block_size = 8           dflash.attention.head_count = 32
dflash.conv_kernel_size = 2     dflash.attention.head_count_kv = 8
dflash.conv_group_size = 16     dflash.attention.key_length = 128
dflash.selector_rank = 256      dflash.selector_top_k = 16
dflash.target_layers = [6, 20, 34, 48, 62]
dflash.attention.sliding_window_pattern = [true, true, true, true, true]

Happy to run a patched build or add instrumentation if that would help narrow it down.

[mirek190]

I am (non-exhaustively and non-scientifically) testing on R9700 + Vulkan.

At spec-draft-n-max = 4 it produced the best results, I see a small bump over MTP (around 49-60 t/s, but averages at around 52 t/s) with some peaks at 74 t/s (fastest I've ever achieved with this model). MTP produces around 40-60t/s but will often sit at around 46-48 t/s.

Those gains are not as spectacular as claimed, but are an improvement. My settings (MTP settings also included, commented out):

./llama-server --jinja --port 8085 -fa on -np 1 --log-colors on --no-webui --metrics --models-max 1 --no-warmup --models-preset models/config.ini

model = AI/LLM/Models/Qwen3.8-27B-UD-Q4_K_XL.gguf
ctx-size = 131072
ctk = iq4_nl
ctv = iq4_nl
ub = 2048
b = 1024
temp = 1.0
top-p = 0.95
top-k = 20
repeat-penalty = 1.0
presence-penalty = 0.0
dev = Vulkan0
chat-template-file=AI/LLM/Models/qwen3.8_chat_template.jinja
#spec-default =
#spec-type = draft-mtp
#spec-draft-n-max = 4
fit = off
spec-type = draft-dflash
spec-draft-n-max = 4
model-draft = AI/LLM/Models/Qwen3.8-27B-DFlash2-Q4_K_M.gguf

The JINJA is froggeric's v22.2

cache q4? LOL

[Zeldatales]

Following up on my earlier Strix Halo numbers with the concurrency data @KaruroChori and @treo asked about. Same machine (Ryzen AI MAX+ 395, gfx1151, Vulkan/RADV), same build (this PR @ 5ecbe1ac1), same Q4_K_M target.

Server started with -np N and --ctx-size 4096*N so per-slot context stays constant. Each cell is 3 rounds of N simultaneous requests, code-generation prompt, max_tokens=300, temperature=0.

Aggregate throughput (tok/s)

Drafter	np=1	np=2	np=4
DFlash2 Q8_0	21.22	15.87	25.32
MTP Q4_0	22.10	16.52	26.35
none	11.11	20.14	34.14

Per-request throughput (tok/s)

Drafter	np=1	np=2	np=4
DFlash2 Q8_0	21.23	8.13	6.53
MTP Q4_0	22.10	8.33	6.80
none	11.11	10.07	8.54

What this shows

1. The baseline scales, the drafted paths don't. Without a drafter, aggregate throughput goes 11.11 → 20.14 → 34.14, close to linear. With either drafter it drops from np=1 to np=2 and only recovers past the single-stream figure at np=4. So the server's batching itself is healthy — the regression is specific to speculative decoding.

2. Speculation is a net loss at np≥2 here. At np=4 both drafters land at 25–26 tok/s aggregate against 34.14 with no drafter at all. The crossover sits between np=1 and np=2. On this hardware speculative decoding is purely a single-stream optimization — it buys 2× there and costs ~25% by np=4.

3. MTP degrades the same way DFlash2 does. @KaruroChori reported that MTP doesn't show this behaviour on Intel B70. On AMD/Vulkan the two are within noise of each other at every concurrency level (np=2: 16.52 vs 15.87; np=4: 26.35 vs 25.32). Whatever causes it here is not specific to DFlash2, so it may be worth checking whether the B70 difference is backend-specific rather than drafter-specific.

Worth noting the failure mode is milder than what was reported on B70 — per-request throughput falls to 6.5–8.3 tok/s rather than collapsing to ~1 tok/s. It degrades rather than falls over.

Caveats

Three rounds per cell, one prompt type, concurrency only up to 4. --cont-batching left at its default (enabled). I didn't vary --spec-draft-n-max across concurrency levels — it stayed at 4 throughout, and a lower value might behave differently under load.

Happy to extend this to np=8, sweep --spec-draft-n-max under concurrency, or test with continuous batching disabled if any of those would be useful.

[ouening]

Tested with folowing enviroment:
OS: Windows 11
CPU：i9 14900kf
RAM: 64G 5400MHz
GPU: RTX 3090 24G
Python: 3.13.12
llama.cpp: pr-27342 (5ecbe1a "support DFlash2")

Baseline
llama-server.exe -m Qwen3.8-27B-UD-Q4_K_M.gguf --fit off -np 1 -fa on -sm none -ts 3,1 -mg 0 -rea off -c 128000 -ctk q8_0 -ctv q8_0 --port 18080

MTP:
llama-server.exe -m Qwen3.8-27B-UD-Q4_K_M.gguf --fit off -np 1 -fa on -sm none -ts 3,1 -mg 0 -rea off -c 128000 -ctk q8_0 -ctv q8_0 --port 18080 --spec-type draft-mtp --spec-draft-n-max (2, 3, 4, 5)

DFlash2
llama-server.exe -m Qwen3.8-27B-UD-Q4_K_M.gguf --fit off -np 1 -fa on -sm none -ts 3,1 -mg 0 -rea off -c 128000 -ctk q8_0 -ctv q8_0 --port 18080 --spec-type draft-dflash --spec-draft-n-max (2, 3, 4, 5) -md Qwen3.8-27B-DFlash2-Q4_K_M.gguf

Test command:
$ python C:\Apps\llama.cpp\tools\server\bench\speed-bench\speed_bench.py --url localhost:8080 --concurrency 1 --limit 2

Results shows bellow:

[take-cheeze]

I've tried this PR with my Ryzen 395+ with claude assistant.
It didn't improve much compared to tuning MTP like https://www.amd.com/en/blogs/2026/run-qwen-3-8-27b-on-amd-ryzen-ai-max-and-radeon-graphics-cards-day-0.html but posting for someones help.

Following is the patch and report written by claude code:

fixdflashmultigpudeviceinherit.patch

DFlash2 on Ryzen AI Max+ 395 (Strix Halo, Radeon 8060S) — test report + multi-GPU crash fix

Tested PR #27342 on a Ryzen AI Max+ 395 / Radeon 8060S ("Strix Halo") box, also equipped with a
discrete RTX 5050 (8GB). Target: Qwen/Qwen3.8-27B Q4_K_M (~15.65 GiB), draft:
incoai/Qwen3.8-27B-DFlash2 Q4_K_M.

Bug: crash on multi-GPU systems when --spec-draft-device isn't set

Running DFlash2 via the Vulkan backend crashed on load whenever the draft model's device
wasn't explicitly pinned to match the target's:

ggml/src/ggml-backend.cpp:930: pre-allocated tensor (output.weight) in a buffer (Vulkan1)
that cannot run the operation (NONE)

(abort in ggml_backend_sched_backend_id_from_cur → graph_reserve → llama_context::llama_context,
happening while building the draft context)

Root cause: DFlash(2) draft models share the target's vocabulary and reuse the target's
output projection tensor directly (src/models/dflash.cpp, output = model_other->output
when the draft has no output weight of its own). That only works if the draft's backend
scheduler actually has the buffer type that tensor lives in — i.e. only if the draft ends up
on the same device as the target.

--device only sets the target's device. The draft's device is the separate
--spec-draft-device/-devd flag, and when it's left unset, common_base_params_to_speculative()
(common/speculative.cpp) unconditionally does:

result.devices = params_spec.devices;   // params_spec.devices is empty when -devd is unset

clobbering the draft's device list — which had already inherited the target's via
common_params result = params; a few lines above — with an empty list, falling back to
whatever the library's default device auto-selection picks. On a single-GPU box that's a
no-op (there's only one device to pick), so the bug is latent there. On this machine, Vulkan
enumerates two devices (Vulkan0 = RTX 5050, Vulkan1 = Radeon 8060S), and auto-selection put
the draft on a different device than the target — hence the crash.

This is pre-existing in shared code, not something the DFlash2 PR introduced — the
clobbering line dates to an earlier commit (f5525f7e7a). DFlash/DFlash2 are just the first
spec types that share weight tensors across the target/draft contexts, so they're the first to
surface it as a hard crash instead of a silent, harmless divergence.

Patch (attached, fix-dflash-multi-gpu-device-inherit.patch): only overwrite the inherited
device list when the draft device was actually specified:

if (!params_spec.devices.empty()) {
    result.devices = params_spec.devices;
}

Verified this fixes the crash — DFlash2 now runs correctly on Vulkan with two GPUs present and
no --spec-draft-device flag needed. Also verified --spec-draft-device still works as an
explicit override when you do want the draft elsewhere.

Performance

Backend	Method	Decode speed	Notes
CUDA (RTX 5050, 8GB)	autoregressive	6.8 t/s	VRAM-starved, spills to CPU
CUDA (RTX 5050, 8GB)	DFlash2	11.6 t/s (~1.7×)
HIP/ROCm (8060S)	autoregressive	11.1 t/s
HIP/ROCm (8060S)	DFlash2	12.2 t/s (~1.1×)
Vulkan/RADV (8060S)	autoregressive	12.7 t/s
Vulkan/RADV (8060S)	DFlash2 (post-fix)	17.6 t/s (~1.4×)

DFlash2's relative speedup is highly hardware-dependent: it helps most when the baseline is
compute/VRAM-starved (RTX 5050 with an 8GB card holding an 18GB model), and least when the
baseline is already fully GPU-resident and bandwidth-efficient (Strix Halo's 128GB unified
memory holds the whole model comfortably, so there's less headroom left to reclaim).

For reference, on the same 8060S/Vulkan setup, MTP (--spec-type draft-mtp, self-speculative,
no separate draft model) reached ~1.5–1.8× over the same autoregressive baseline — noticeably
ahead of DFlash2 here, likely because MTP's draft head shares almost the entire trunk with the
target (near-free verification), whereas DFlash2's separate draft network pays a full second
forward pass every block. Worth noting as a data point for where DFlash2 is and isn't the
right tool, not a criticism of the PR — different architectures trading off differently across
hardware is expected.

Environment

• CPU/APU: AMD Ryzen AI Max+ 395 (Strix Halo), Radeon 8060S Graphics (gfx1151), 128GB GTT
• Discrete GPU: NVIDIA GeForce RTX 5050, 8GB VRAM
• Backends tested: CUDA 13.3, ROCm/HIP 7.14, Vulkan (RADV, mesa 26.1.4, glslc 2026.1)
• Base commit: 5ecbe1ac1 (PR spec : add DFlash2 support (local convolution + candidate selector) #27342 head at time of testing)

[Stoney49th]

I've tried this PR with my Ryzen 395+ with claude assistant. It didn't improve much compared to tuning MTP like https://www.amd.com/en/blogs/2026/run-qwen-3-8-27b-on-amd-ryzen-ai-max-and-radeon-graphics-cards-day-0.html but posting for someones help.

Following is the patch and report written by claude code:

fixdflashmultigpudeviceinherit.patch

DFlash2 on Ryzen AI Max+ 395 (Strix Halo, Radeon 8060S) — test report + multi-GPU crash fix

Tested PR #27342 on a Ryzen AI Max+ 395 / Radeon 8060S ("Strix Halo") box, also equipped with a discrete RTX 5050 (8GB). Target: Qwen/Qwen3.8-27B Q4_K_M (~15.65 GiB), draft: incoai/Qwen3.8-27B-DFlash2 Q4_K_M.

Bug: crash on multi-GPU systems when --spec-draft-device isn't set

Running DFlash2 via the Vulkan backend crashed on load whenever the draft model's device wasn't explicitly pinned to match the target's:

ggml/src/ggml-backend.cpp:930: pre-allocated tensor (output.weight) in a buffer (Vulkan1)
that cannot run the operation (NONE)

(abort in ggml_backend_sched_backend_id_from_cur → graph_reserve → llama_context::llama_context, happening while building the draft context)

Root cause: DFlash(2) draft models share the target's vocabulary and reuse the target's output projection tensor directly (src/models/dflash.cpp, output = model_other->output when the draft has no output weight of its own). That only works if the draft's backend scheduler actually has the buffer type that tensor lives in — i.e. only if the draft ends up on the same device as the target.

--device only sets the target's device. The draft's device is the separate --spec-draft-device/-devd flag, and when it's left unset, common_base_params_to_speculative() (common/speculative.cpp) unconditionally does:

result.devices = params_spec.devices;   // params_spec.devices is empty when -devd is unset

clobbering the draft's device list — which had already inherited the target's via common_params result = params; a few lines above — with an empty list, falling back to whatever the library's default device auto-selection picks. On a single-GPU box that's a no-op (there's only one device to pick), so the bug is latent there. On this machine, Vulkan enumerates two devices (Vulkan0 = RTX 5050, Vulkan1 = Radeon 8060S), and auto-selection put the draft on a different device than the target — hence the crash.

This is pre-existing in shared code, not something the DFlash2 PR introduced — the clobbering line dates to an earlier commit (f5525f7e7a). DFlash/DFlash2 are just the first spec types that share weight tensors across the target/draft contexts, so they're the first to surface it as a hard crash instead of a silent, harmless divergence.

Patch (attached, fix-dflash-multi-gpu-device-inherit.patch): only overwrite the inherited device list when the draft device was actually specified:

if (!params_spec.devices.empty()) {
    result.devices = params_spec.devices;
}

Verified this fixes the crash — DFlash2 now runs correctly on Vulkan with two GPUs present and no --spec-draft-device flag needed. Also verified --spec-draft-device still works as an explicit override when you do want the draft elsewhere.

Performance

Backend Method Decode speed Notes
CUDA (RTX 5050, 8GB) autoregressive 6.8 t/s VRAM-starved, spills to CPU
CUDA (RTX 5050, 8GB) DFlash2 11.6 t/s (~1.7×)
HIP/ROCm (8060S) autoregressive 11.1 t/s
HIP/ROCm (8060S) DFlash2 12.2 t/s (~1.1×)
Vulkan/RADV (8060S) autoregressive 12.7 t/s
Vulkan/RADV (8060S) DFlash2 (post-fix) 17.6 t/s (~1.4×)

DFlash2's relative speedup is highly hardware-dependent: it helps most when the baseline is compute/VRAM-starved (RTX 5050 with an 8GB card holding an 18GB model), and least when the baseline is already fully GPU-resident and bandwidth-efficient (Strix Halo's 128GB unified memory holds the whole model comfortably, so there's less headroom left to reclaim).

For reference, on the same 8060S/Vulkan setup, MTP (--spec-type draft-mtp, self-speculative, no separate draft model) reached ~1.5–1.8× over the same autoregressive baseline — noticeably ahead of DFlash2 here, likely because MTP's draft head shares almost the entire trunk with the target (near-free verification), whereas DFlash2's separate draft network pays a full second forward pass every block. Worth noting as a data point for where DFlash2 is and isn't the right tool, not a criticism of the PR — different architectures trading off differently across hardware is expected.

Environment

* CPU/APU: AMD Ryzen AI Max+ 395 (Strix Halo), Radeon 8060S Graphics (gfx1151), 128GB GTT

* Discrete GPU: NVIDIA GeForce RTX 5050, 8GB VRAM

* Backends tested: CUDA 13.3, ROCm/HIP 7.14, Vulkan (RADV, mesa 26.1.4, glslc 2026.1)

* Base commit: `5ecbe1ac1` (PR [spec : add DFlash2 support (local convolution + candidate selector) #27342](https://github.com/ggml-org/llama.cpp/pull/27342) head at time of testing)

I see the same scaling issue. Slightly elevated performance compared to MTP on 2xR9700 @ Vulkan and ROCM (50/50 Split) of Q3.8-27B-UD-Q4-XL. Maybe 10-15%, and I also noticed a crazy spike to almost 4x when all of the predictions matched. Well done :)

However, with parallel sessions, I see exactly the same steep drop as with MTP with vulkan and ROCM. Which is really cumbersome for any real-world workload. But thats not subject to the PR i guess but a general bug right now somewhere in the server or seculation implementation.

I did not test the mtmd path.

[ix-yuzhoulu]

Incompatibility between draft-dflash and Tensor Parallelism (-sm tensor)

Description

Using -sm tensor together with --spec-type draft-dflash triggers an assertion failure during the draft model loading stage:

ggml/src/ggml-backend-meta.cpp:543: GGML_ASSERT(src_ss[0].axis != GGML_BACKEND_SPLIT_AXIS_0) failed

The stack trace indicates that in ggml_backend_meta_get_split_state, the tensor's split axis is unexpectedly set to 0 (GGML_BACKEND_SPLIT_AXIS_0).

Impact & Workaround

Currently, switching to -sm layer serves as a workaround. However, pipeline bubbles introduced by layer-wise splitting cause generation throughput to drop significantly from the expected **80+ t/s ** (Measured by MTP) down to ~18 t/s.

Reproduction Command

llama-server \
  -m Qwen3.8-27B-Q4_K_M.gguf \
  --model-draft Qwen3.8-27B-DFlash2-Q4_K_M.gguf \
  --spec-type draft-dflash \
  --spec-draft-n-max 7 \
  -ngl 99 -sm tensor -c 262144 \
  --parallel 2 -fa on -ctk q8_0 -ctv q8_0 --kv-unified

[gordan-bobic]

This errors out with more than 1 GPU and tensor parallelism. Tried 2x 3x 4x 5x GPUs:

/var/lib/ollama/llama.cpp/ggml/src/ggml-backend-meta.cpp:575: GGML_ASSERT(src_ss[0].axis != GGML_BACKEND_SPLIT_AXIS_0) failed

[lexasub]

relates to #27396

[savageops]

A patch set for review: ngram cache overlay, loader bug, p_min wiring (measured notes)

First — thank you for this PR. DFlash2 on Qwen3.8-27B is excellent
work; the selector-lattice design drafts noticeably better than the
built-in MTP on our rig (acceptance ~4.5-5.5 vs ~2.3 mean).

We've been running it hard on a single RTX 5090 (Qwen3.8-27B Q5_K_M
target, DFlash2 Q4_K_M sidecar, 131K ctx) and built a small set of
additions on top. Everything below is measured with paired
interleaved A/B rounds (sequential-statistics stop, verdicts with
CIs) and byte-identity greedy gates. Full measured record, including everything that did NOT work:
https://gist.github.com/savageops/c29d8048a0a46c1e60c9566a8d87c8c8 Happy to split
into separate patches or rework any of this.

1. An ngram cache overlay inside the dflash impl (the main addition)

draft-dflash,ngram-cache currently drops the ngram impl silently.
Rather than touching the impl-selection layer, we embed a
context-keyed ngram cache inside the dflash drafter (it already
receives the full live context in dp.prompt):

• chain >= 8 tokens -> draft 16-wide, skip the 2B block decode
• chain 4-7 -> draft capped at 8
• anything less -> the lattice as normal

Rationale (measured on GPU-resident targets): verify batches are
~flat in length, so long confident drafts are cheap and long
weak-signal drafts are not. Gate on chain confidence, never on
per-token probability.

Config: requires --spec-draft-n-max 16 (the framework sizes verify
buffers from n_max; we pin the lattice loop to the checkpoint's
trained n_extract in code so the drafter still emits width-5).

workload (wall tok/s)	dflash @ n7	+ overlay @ n16
exact-repeat content	~95	216-224
openapi/yaml schema gen	82.6	155-161
json object arrays	137	150-167
python class w/ methods	111	125-138
4-workload general avg	112	118

Greedy outputs byte-identical on our merge probe (one caveat below).

2. Bug: n_max=4 crashes the draft loader

Deterministic repro: --spec-draft-n-max 4 with the current GGUF ->
done_getting_tensors: wrong number of tensors; expected 81, got 58
(the draft GGUF gets resolved against the target's tensor map under
some load paths). n_max 5/6/7/16 load fine. Happy to file separately
with logs if useful.

3. Smaller items

• The DFlash2 lattice branch ignores --spec-draft-p-min (the
  DFlash1/DSpark branches honor it). We wired it through; note that
  on GPU-resident targets confidence truncation measured strictly
  WORSE for us (explain 76 -> 45 tok/s — verify batches are flat, so
  wasted tail tokens cost ~nothing while truncated good drafts force
  more passes). Including it as an option for CPU-bound lanes only.
• n_max == trained extract width is the lattice's performance peak
  (curve: 5 peak / 6 erratic / 7 -5% / 4 crash). Might be worth a
  log line suggesting it.
• Never quantize the drafter's KV: --spec-draft-type-k q8_0 cost us
  76 -> 51 tok/s post-100K prefill. The 2B's F16 KV precision is
  load-bearing at depth.

4. One honest caveat + what we could NOT do

At draft width > trained width, the wider verify batch changes CUDA
graph shape and near-tie argmaxes can flip — we observed a single
reordered line in one probe (semantically identical). Byte-identity
is a property of the narrow config.

We also attempted cross-stream pre-drafting (worker thread decoding
the next block into ctx_dft during target verify). Protocol worked
(positions/KV handled), but concurrent llama_decode from two host
threads fails on the shared CUDA backend, and on a CPU-backend
drafter it crashes the server. We measured the pipelining upside at
+10-45% on our rig (target-only 62.4 tok/s baseline; pass
decomposition ~18ms verify + ~8ms draft) — if the backend ever gains
per-context threadpool isolation or multi-stream support, that
ceiling is real and the spec-decode layer should be ready for it.

---

Repro commands and the full falsification table (truncation, EWMA
gate controller - measured neutral, ngram width sweep, KV quant,
suffix-automaton drafter) are in the writeup: https://gist.github.com/savageops/c29d8048a0a46c1e60c9566a8d87c8c8 Patch commits are
rebased on 5ecbe1a. Thanks again — happy to restructure any of
this to match how you'd like it upstream.

[ramlok0]

llama.cpp crashed when testing this branch, not sure if it is related to your changes or useful (core was not produced), just fyi.

build done from commit 5ecbe1a
command line
llama-server --temp 1 --top-p 0.95 --min-p 0.00 --top-k 40 --jinja --repeat-penalty 1.05 --presence-penalty 0.0 -np 1 --threads 10 --threads-batch 10 --spec-type draft-dflash --spec-draft-n-max 10 --chat-template-file chat_template_sharp.jinja --batch-size 4096 --ubatch-size 4096 -m Qwen3.8-27B-UD-Q8_K_XL.gguf -md Qwen3.8-27B-DFlash2-Q8_0.gguf -ngl 999 -fa on --load-mode dio -c 240000 --reasoning-preserve --reasoning auto --chat-template-kwargs {"preserve_thinking":"true","reasoning_effort":"xhigh"}

ROCm 7.14, Ubuntu 24.04, 2x W7800 48GB cards

74.17.038.230 E ROCm error: an illegal memory access was encountered
/home/pc/tools/llama.cpp_rocm_dflash2/ggml/src/ggml-cuda/ggml-cuda.cu:106: ROCm error
74.17.038.236 E   current device: -1, in function ggml_backend_cuda_buffer_set_tensor at /home/pc/tools/llama.cpp_rocm_dflash2/ggml/src/ggml-cuda/ggml-cuda.cu:786
74.17.038.236 E   hipMemcpyAsync((char *) tensor->data + offset, data, size, hipMemcpyHostToDevice, ((hipStream_t)2))


This GDB supports auto-downloading debuginfo from the following URLs:
  <https://debuginfod.ubuntu.com>
Enable debuginfod for this session? (y or [n]) [answered N; input not from terminal]
Debuginfod has been disabled.
To make this setting permanent, add 'set debuginfod enabled off' to .gdbinit.
[Thread debugging using libthread_db enabled]
Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
0x000070f59e510913 in __GI___wait4 (pid=16854, stat_loc=0x0, options=0, usage=0x0) at ../sysdeps/unix/sysv/linux/wait4.c:30
warning: 30	../sysdeps/unix/sysv/linux/wait4.c: No such file or directory
#0  0x000070f59e510913 in __GI___wait4 (pid=16854, stat_loc=0x0, options=0, usage=0x0) at ../sysdeps/unix/sysv/linux/wait4.c:30
30	in ../sysdeps/unix/sysv/linux/wait4.c
#1  0x000070f59eb1f6c3 in ggml_print_backtrace () from /home/pc/tools/llama.cpp_rocm_dflash2/build/bin/libggml-base.so.0
#2  0x000070f59eb1f86b in ggml_abort () from /home/pc/tools/llama.cpp_rocm_dflash2/build/bin/libggml-base.so.0
#3  0x000070f59c67ee02 in ggml_cuda_error(char const*, char const*, char const*, int, char const*) () from /home/pc/tools/llama.cpp_rocm_dflash2/build/bin/libggml-hip.so.0
#4  0x000070f59c682057 in ggml_backend_cuda_buffer_set_tensor(ggml_backend_buffer*, ggml_tensor*, void const*, unsigned long, unsigned long) () from /home/pc/tools/llama.cpp_rocm_dflash2/build/bin/libggml-hip.so.0
#5  0x000070f59db087d9 in llama_io_read_host::~llama_io_read_host() () from /home/pc/tools/llama.cpp_rocm_dflash2/build/bin/libllama.so.0
#6  0x000070f59dafc16b in llama_context::state_seq_set_data(int, unsigned char const*, unsigned long, unsigned int) () from /home/pc/tools/llama.cpp_rocm_dflash2/build/bin/libllama.so.0
#7  0x000070f59e033e7d in common_prompt_checkpoint::load_dft(llama_context*, int, unsigned int) const () from /home/pc/tools/llama.cpp_rocm_dflash2/build/bin/libllama-common.so.0
#8  0x000070f59edc2640 in server_context_impl::pre_decode()::{lambda(server_slot&)#6}::operator()(server_slot&) const () from /home/pc/tools/llama.cpp_rocm_dflash2/build/bin/libllama-server-impl.so
#9  0x000070f59edb6137 in server_context_impl::iterate(std::vector<server_slot, std::allocator<server_slot> >&, std::function<void (server_slot&)>) () from /home/pc/tools/llama.cpp_rocm_dflash2/build/bin/libllama-server-impl.so
#10 0x000070f59edb6787 in server_context_impl::pre_decode() () from /home/pc/tools/llama.cpp_rocm_dflash2/build/bin/libllama-server-impl.so
#11 0x000070f59edba07c in server_context_impl::update_slots() () from /home/pc/tools/llama.cpp_rocm_dflash2/build/bin/libllama-server-impl.so
#12 0x000070f59ed6124d in server_queue::start_loop(long) () from /home/pc/tools/llama.cpp_rocm_dflash2/build/bin/libllama-server-impl.so
#13 0x000070f59ecfed33 in llama_server(common_params&, int, char**) () from /home/pc/tools/llama.cpp_rocm_dflash2/build/bin/libllama-server-impl.so
#14 0x000070f59ecffd56 in llama_server(int, char**) () from /home/pc/tools/llama.cpp_rocm_dflash2/build/bin/libllama-server-impl.so
#15 0x000070f59e42a1ca in __libc_start_call_main (main=main@entry=0x654f423e9270 <main>, argc=argc@entry=59, argv=argv@entry=0x7ffe7fbb4f18) at ../sysdeps/nptl/libc_start_call_main.h:58
warning: 58	../sysdeps/nptl/libc_start_call_main.h: No such file or directory
#16 0x000070f59e42a28b in __libc_start_main_impl (main=0x654f423e9270 <main>, argc=59, argv=0x7ffe7fbb4f18, init=<optimized out>, fini=<optimized out>, rtld_fini=<optimized out>, stack_end=0x7ffe7fbb4f08) at ../csu/libc-start.c:360
warning: 360	../csu/libc-start.c: No such file or directory
#17 0x0000654f423e92a5 in _start ()
[Inferior 1 (process 12706) detached]
 12706 Aborted                 (core dumped) "${CMD_ARGS[@]}"

I have not had any crashes running latest master with mtp/ngram.

[chuanyu926]

Root cause for the invalid vector subscript draft-load failure that @ThConMan reported above, plus a fix. Their diagnosis that it is an .at() overrun rather than an operator[] was correct and is what pointed me at the right family of call sites.

TL;DR

splits[i] /= split_sum in llama_model_base::load_tensors (src/llama-model.cpp:1351) is not guarded against split_sum == 0. A draft model is loaded after the target has already filled VRAM, so CUDA reports 0 bytes free, splits[0] becomes 0/0 = NaN, every upper_bound comparison against NaN is false, and the returned index runs one past the last device at src/llama-model.cpp:1363.

It is a VRAM-pressure bug, not a checkpoint bug — which is why it reproduces for some people and not others with the identical GGUF.

Environment

build 10498 (5ecbe1ac1), this PR, clean checkout + instrumentation
Windows 11, MSVC 19.41 (VS 2022 BuildTools), CUDA 12.8
cmake -B build -G Ninja -DCMAKE_BUILD_TYPE=Release -DGGML_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=89
RTX 4070 Ti SUPER (sm_89, 16 GB)
target: ggml-org/Qwen3.8-27B-GGUF Q4_K_M
drafter: incoai/Qwen3.8-27B-DFlash2-GGUF Q4_K_M

Narrowing it down

Phase markers inside llama_model_load put the throw inside model->load_tensors(ml), ~80 ms after entry — i.e. during tensor creation, before any tensor data is read. Note the exception message has no prefix, which already rules out load_hparams and load_vocab since both wrap their own exceptions.

Adding a bounds check in front of devices.at(layer_gpu) gives:

DBG devices OOR: layer_gpu=1 devices.size=1 il=0
llama_model_load: error loading model: invalid vector subscript

One device, and upper_bound returned index 1.

The mechanism

// src/llama-model.cpp, load_tensors
splits[i] = free;                    // 1338 - free VRAM in bytes

float split_sum = 0.0f;
for (size_t i = 0; i < n_devices(); ++i) {
    split_sum += splits[i];
    splits[i] = split_sum;
}
for (size_t i = 0; i < n_devices(); ++i) {
    splits[i] /= split_sum;          // 1351 - 0/0 -> NaN when free == 0
}
...
const int layer_gpu = std::upper_bound(splits.begin(), splits.begin() + n_devices(),
                          float(il - i_gpu_start)/act_gpu_layers) - splits.begin();
auto * dev = devices.at(layer_gpu).dev;   // 1363 - throws

With a single device and free == 0: splits[0] = 0/0 = NaN. upper_bound looks for the first element where value < element; 0.0 < NaN is false, so it walks off the end and yields n_devices(). .at() then throws std::out_of_range, which MSVC spells invalid vector subscript.

The existing fallback just above does not catch this:

if (free == 0 && total == 0) {       // 1335
    ggml_backend_dev_memory(cpu_dev, &free, &total);
}

total is the full 16 GB here, so the condition is false and splits[0] = 0 stands.

This is reachable by any second model loaded onto a saturated device, but speculative decoding is the case that hits it routinely, since the drafter loads last by construction.

Fix

-    for (size_t i = 0; i < n_devices(); ++i) {
-        splits[i] /= split_sum;
+    if (split_sum > 0.0f) {
+        for (size_t i = 0; i < n_devices(); ++i) {
+            splits[i] /= split_sum;
+        }
+    } else {
+        // every device reported zero free memory - this happens when a draft model is
+        // loaded after the target has already filled VRAM. dividing by zero here would
+        // poison the split points with NaN and send upper_bound() past the last device.
+        for (size_t i = 0; i < n_devices(); ++i) {
+            splits[i] = float(i + 1) / float(n_devices());
+        }
     }

and defensively at 1363:

-        auto * dev = devices.at(layer_gpu).dev;
+        auto * dev = devices.at(std::min<size_t>(layer_gpu, n_devices() - 1)).dev;

Either half alone stops the crash; the first is the actual bug, the second keeps a NaN from any other source from turning into an out-of-range access. Happy to open this as a separate PR if you prefer it split out from the DFlash2 work, since it is really a pre-existing loader issue that DFlash2 only exposes.

After the fix

Loads and runs. -ngl 32 of 65 layers on GPU, -ngld 99, -c 8192 -fa on -ctk q8_0 -ctv q8_0 --fit off -np 1, --spec-draft-n-max 5 (matching the checkpoint's n_extract=5), greedy, same prompts run against the same build with and without -md:

workload	target only	+ DFlash2	speedup	draft acceptance
code generation	4.05 tok/s	8.23 tok/s	2.03x	65.2%
prose explanation	3.96 tok/s	4.92 tok/s	1.24x	29.4%

Absolute numbers are low because half the target sits on CPU on a 16 GB card — worth noting that speculative decoding holds up well under partial offload, since the verify pass batches tokens and amortizes the slow layers. The relative gain on code is larger than what I get on a fully GPU-resident setup would suggest is the ceiling.

Unrelated observation while tracing this

common/speculative.cpp:1031 (DFlash) calls

llama_set_embeddings_layer_inp(ctx_tgt, (uint32_t) target_layer_ids[k], true);

with no bounds check, while the EAGLE3 path at line 516 guards the same call with target_layer_ids[k] < n_layer_tgt and aborts with a clear message otherwise. Not something I hit — the Qwen3.8-27B drafter's taps [6, 20, 34, 48, 62] are all in range — but a drafter whose target_layers exceed the target's depth would fail there without a useful diagnostic.

[Stoney49th]

Incompatibility between draft-dflash and Tensor Parallelism (-sm tensor)

Description

Using -sm tensor together with --spec-type draft-dflash triggers an assertion failure during the draft model loading stage:

ggml/src/ggml-backend-meta.cpp:543: GGML_ASSERT(src_ss[0].axis != GGML_BACKEND_SPLIT_AXIS_0) failed

The stack trace indicates that in ggml_backend_meta_get_split_state, the tensor's split axis is unexpectedly set to 0 (GGML_BACKEND_SPLIT_AXIS_0).

Impact & Workaround

Currently, switching to -sm layer serves as a workaround. However, pipeline bubbles introduced by layer-wise splitting cause generation throughput to drop significantly from the expected **80+ t/s ** (Measured by MTP) down to ~18 t/s.

Reproduction Command

llama-server \
  -m Qwen3.8-27B-Q4_K_M.gguf \
  --model-draft Qwen3.8-27B-DFlash2-Q4_K_M.gguf \
  --spec-type draft-dflash \
  --spec-draft-n-max 7 \
  -ngl 99 -sm tensor -c 262144 \
  --parallel 2 -fa on -ctk q8_0 -ctv q8_0 --kv-unified

@ix-yuzhoulu can you elaborate more on these pipeline bubbles? I ried to narrow it down myself but only got so far that is somehow tied to the number of extra tokens generated in relation to the number of slots, thats why currently MTP with draft n max = 1 scales and others dont. I saw the same here with Dflash2. But it seems like CUDA is not affected at all, while vulkan is badly struck and ROCM kind of - at least with MTP, i have subpar scaling. it scales however somewhat, and funny enough, for longer token sequences ,it becomes faster. I could not test due to the model loading crash, but I'm willing to retest after the fix. Heres the scaling for ROCM with MTP and vulkan, current master, for MTP. so I expect it to be roughly the same for dflash2:

Details

The external CUDA datapoint in the comments supplies the comparison this section previously lacked. Placing
the two cost curves side by side, at comparable widths:

width	Vulkan (2× R9700)	ROCm (2× R9700, §13)	CUDA (RTX PRO 4000, 70 W)
2	49 ms	52.8 ms	51.4 ms
4	63 / 78 ms	76.8 / 78.3 ms	75.9 ms
6	89.3 ms	99.3 ms	—
8	110 / 120 ms	122.0 ms	—
9	162.8 ms	— (n_max=2 not yet run)	—
12	170 ms	93.4 ms ⚠ unreproduced	~79 ms
32	—	—	~143 ms

ROCM

Run R2 — spec-draft-n-max = 3 (L = 4), parallel = 3

R2a — single (slot 2, task 0)

metric	value
prompt eval	300.74 ms / 14 tok — 21.48 ms/tok — 46.55 t/s
eval	21648.87 ms / 712 tok — 30.45 ms/tok — 32.84 t/s
graphs reused	276
draft acceptance	0.52278 (436/834), mean len 2.57
acc per pos	(0.719, 0.496, 0.353)

R2b — 2 parallel (tasks 281, 282)

slot	task	eval	ms/tok	t/s	graphs reused	acceptance	mean len
2	281	22204.66 ms / 374 tok	59.53	16.80	445	0.393 (203/516)	2.18
1	282	27441.32 ms / 505 tok	54.45	18.37	514	0.365 (265/726)	2.10

R2c — 3 parallel (tasks 527, 528, 529)

slot	task	eval	ms/tok	t/s	graphs reused	acceptance	mean len
0	529	17070.45 ms / 455 tok	37.60	26.60	703	0.462 (265/573)	2.39
2	528	20411.33 ms / 461 tok	44.37	22.54	726	0.381 (247/648)	2.14
1	527	20822.56 ms / 465 tok	44.88	22.28	730	0.367 (243/663)	2.10

Aggregate TG: single 32.84 → dual 35.17 (1.07×) → triple 71.42 (2.18×).
Tick time: width 4 = 78.3 ms · width 8 = 122.0 ms · width 12 = 93.4 ms.

Width 12 costing less than width 8 is the opposite of Vulkan's width-12 point (Vulkan Run 2, same
config: single 37.52, triple 40.78, ratio 1.09×, tick time 170.0 ms at width 12 — rising, not
falling, from its own width-8-adjacent points). Task 529 (slot 0) alone ran at 37.60 ms/tok, well
below the other two triple-leg slots (44.4–44.9 ms/tok) despite all three submitting into the same
verify batch each tick — a per-slot spread not seen in the Vulkan runs at this width.

Run R1 — spec-draft-n-max = 1 (L = 2), parallel = 3

R1a — single (slot 2, task 0)

metric	value
prompt eval	357.63 ms / 14 tok — 25.55 ms/tok — 39.15 t/s
eval	17343.49 ms / 563 tok — 30.86 ms/tok — 32.40 t/s
graphs reused	327
draft acceptance	0.70821 (233/329), mean len 1.71
acc per pos	(0.708)

R1b — 2 parallel (tasks 332, 333)

slot	task	eval	ms/tok	t/s	graphs reused	acceptance	mean len
2	333	21937.01 ms / 498 tok	44.14	22.66	606	0.762 (215/282)	1.76
1	332	22918.86 ms / 503 tok	45.66	21.90	626	0.657 (199/303)	1.66

R1c — 3 parallel (tasks 639, 640, 641)

slot	task	eval	ms/tok	t/s	graphs reused	acceptance	mean len
0	641	23564.45 ms / 400 tok	59.06	16.93	860	0.691 (163/236)	1.69
2	639	23807.62 ms / 393 tok	60.73	16.47	860	0.658 (156/237)	1.66
1	640	24874.73 ms / 426 tok	58.53	17.09	879	0.658 (169/257)	1.66

Aggregate TG: single 32.40 → dual 44.56 (1.38×) → triple 50.49 (1.56×).
Tick time (ms/tok × mean_len): width 2 = 52.8 ms · width 4 = 76.8 ms · width 6 = 99.3 ms.

ulkan comparison at the same config (Runs 1/4): single 34–35 t/s, triple 52.7–55.8 t/s (1.50–1.64×),
tick time 49.7 ms (width 2) → 85.5–89.3 ms (width 6). **ROCm tracks the Vulkan collapse pattern

Vulkan

Sweep summary (raw)

run	n_max	L	N	N·L	single TG	parallel TG (sum)	parallel/single	single acc	parallel acc (range)
6a	1	2	2	4	33.72	42.96	1.27×	0.644	0.661 – 0.686
1	1	2	3	6	35.11	52.65	1.50×	0.699	0.656 – 0.780
4	1	2	3	6	34.07	55.78	1.64×	0.693	0.643 – 0.675
6b	3	4	2	8	35.42	41.12	1.16×	0.416	0.415 – 0.425
5	1	2	4	8	33.87	55.58	1.64×	0.632	0.632 – 0.683
3	2	3	3	9	35.95	38.52	1.07×	0.529	0.530 – 0.560
2	3	4	3	12	37.52	40.78	1.09×	0.458	0.393 – 0.500

Runs 6b and 5 are the same batch width (8 tokens) with different L — 4 vs 2 — and differ by
1.16× vs 1.64×. Recorded here without interpretation.

Config

[qwen3-8-27b]
cache-ram = 14336
ctx-checkpoints = 4
checkpoint-min-step = 8192
main-gpu = 0
parallel = 3
batch-size = 4096
ubatch-size = 512
kv-unified = false
spec-type = draft-mtp
spec-draft-n-max = 1
hf = unsloth/Qwen3.8-27B-GGUF:UD-Q4_K_XL
ctx-size = 491520
temp = 1.0
top-p = 0.95
top-k = 20
min-p = 0.0
presence-penalty = 0.0
repeat-penalty = 1.0
cache-type-k = q8_0
cache-type-v = q8_0
flash-attn = true
split-mode = layer
tensor-split = 50,50,0
jinja = true
reasoning-preserve = true
reasoning-effort = medium
reasoning-budget = 20000
image-min-tokens = 1024
no-mmproj-offload = true