cuda: Bonsai 2 27B at the full 262k window with the MTP head on 12 GB Ada: #218 + #215 hybrid PTQ1_0 dispatch, in-place q4_0/q8_0 K/V flash attention (integration checkpoint)

[professorpalmer]

Summary

Integration checkpoint: Bonsai 2 27B served at its full 262,144-token window with the MTP draft head, on a 12 GB RTX 4070, at 103 tok/s greedy (three-prompt mean, up from 60 without the draft head at this context and from 51 on the stock build; 121 on code). Same PTQ1_0 weights, q4_0 K/V, byte-identical output.

It is the merge of two independent PTQ1_0 lines of work plus two new pieces that neither had:

commit	author	what
#217 31ad620	@sudoingX	Hadamard inverse on the MTP token embeddings (no unrotated embedding copy)
#218 98ea410..5883186	@sudoingX	planar-transposed activations, dedicated multi-column PTQ1_0 mat-vec, GGML_CUDA_BATCH_INVARIANT, bf16 small-N mat-vec
#220 82ebc3d	me	recurrent-state gather folded into GATED_DELTA_NET, per-context registry
#214 94c913f	me	branch-free PTQ1_0 MMQ tile loader, full Ampere tile table (2x prefill)
#216 e6b40a9	me	4-column GDN warp layout on all Ampere+
#215 6f856a7	me	SoA q8 activations + warp-per-row small-K GEMV for the one-column decode path, rebased onto #218
new 29ec099	me	hybrid dispatch: one layout enum decides per call; PTQ1_0 mat-vec up to 4 columns, MMQ from 5
new e5cd0f3	me	flash attention MMA reads q4_0/q8_0 K/V in place (no F16 scratch conversion)
new 847a323	me	multi-column PTQ1_0 mat-vec: raw digits with exact activation sums, per-pair epilogue (+13% 3-col, +24% 4-col)
new 1e1409a	me	qwen35 MTP graph: publish only the output rows of h_nextn when masked (latent bug, see below)
new d8a7180	me	draft-mtp: catch-up rows decoded together with the first draft row (one launch and one host round trip fewer per step)

His commits are cherry-picked with authorship intact. The standalone PRs stay open for review on their own; this branch is what all of them look like together and is what produced the numbers below. Reviewers can take it whole or use it to pick pieces.

The new commits

Hybrid PTQ1_0 dispatch (29ec099)

Measured head to head on this card, the #215 layout wins at one column and the #218 kernel wins at 2-4 columns, so both are kept and one function chooses:

enum ggml_cuda_q8_1_layout { GGML_CUDA_Q8_1_AOS, GGML_CUDA_Q8_1_SOA_ISUM, GGML_CUDA_Q8_1_PT };
// PTQ1_0: one column, no ids -> SOA_ISUM (decode hot path); 2-8 columns or MoE ids -> PT.
// GGML_CUDA_BATCH_INVARIANT=1 forces one column onto PT so it runs the same arithmetic as the verify pass.

Quantizer and kernel both call ggml_cuda_q8_1_layout_host(), so they cannot disagree. y_soa is a template parameter of mul_mat_vec_q (a runtime flag cost registers and 4% of decode). ggml_cuda_should_use_mmvq routes PTQ1_0 to the mat-vec path through ne11 <= 4; the MMQ tile path (#214) takes over from 5 columns. calc_nwarps/calc_rows_per_block apply the #218 4-warp/1-row geometry only for ncols_dst > 1.

Flash attention on quantized K/V without the F16 copy (e5cd0f3)

With -ctk q4_0 -ctv q4_0 the MMA kernel used to convert the whole K/V cache to F16 into a scratch buffer per micro-batch during prefill. At 262k that is ~1 GB per model, ~2 GB with the MTP draft, which is exactly what did not fit in 12 GB. The MMA kernel now takes type_KV as a template parameter: flash_attn_ext_f16_load_tile_q dequantizes q4_0/q8_0 blocks straight from the cache into the shared-memory tile (aligned 32-bit word reads, scale + nibble/byte unpack to half2), the cp.async multi-stage pipeline is disabled for quantized types (it cannot dequantize), strides become byte strides, and ggml_cuda_flash_attn_ext_get_alloc_size no longer reserves the F16 buffers when the in-place kernel applies. Instances are generated for D=128 and D=256, ncols 8..64, q4_0 and q8_0. ggml_cuda_fattn_mma_kv_native_supported() gates it on head size, matching K/V types and word-aligned strides; everything else takes the old path unchanged.

Multi-column PTQ1_0 mat-vec, second pass (847a323)

Nsight Compute on the 3-column verify kernel from #218 showed it latency-bound at 25% occupancy, 5.6 warp cycles per issued instruction and ~2100 dynamic instructions per warp where ~1300 do the dot products. Two changes, arithmetic unchanged:

• ptq1_0_trit_step returns the raw 0/1/2 digits instead of biasing them to -1/0/1 per word. The dot is dp4a(unsigned digits, signed q8) and the bias is applied once per 32-block as sumi - isum, where isum is the exact int16 sum of the quantized activations that quantize_q8_1 now stores in the ds.y half slot for the PT layout (bit-cast, same as the SOA_ISUM layout does). One SIMD subtract per 4 weights gone.
• The epilogue was one warp per (row, column) doing a lane-strided sum plus a shuffle butterfly, with a runtime division for the row index. It is now one thread per (row, column) pair summing the K-block partials from shared memory with four interleaved accumulators, a padded partials stride (bpr + 1) so there are no bank conflicts, and a fastdiv for rows_per_cta.

Instructions per warp ~2100 -> ~1500, DRAM utilisation 61% -> 83%, stall samples on the epilogue shuffles gone. llama-bench on this branch at the serving memory clock: pp3 145.4 -> 169.5, pp4 167.1 -> 212.2, pp1 unchanged (different kernel). Tried and reverted with numbers in the linked repo: a cp.async shared-memory stage for the weights (-7%), a 128-register cap for 4 CTAs/SM (spills, -5%), and 2 rows per thread at double occupancy (-19%); the 4-row, 3-CTA geometry stays.

MTP h_nextn rows (1e1409a) and the merged catch-up (d8a7180)

Per step the draft side was four launches: the target verify graph, a catch-up decode of the draft over every verified row (to fill the draft's memory; partly on rows the target rejects), then two single-row draft passes. A CUPTI trace of a step showed each launch boundary costing 0.1-0.45 ms of host time on WDDM on top of the draft's ~0.85 ms GPU time for the catch-up. d8a7180 keeps the catch-up rows in process() and decodes only the accepted prefix at the start of draft() in the same llama_decode as the first draft row. Positions are contiguous and the draft row attends to them in-batch, so the drafts are the same tokens; acceptance per prompt and step counts are identical to the eager path, and output stays byte-identical to the draft head off. Prompts fed in tiny ubatches, a sequence that stops drafting, and rollback replay fall back to a separate decode. Shared-memory (gemma4) and chained-head drafters keep the old path; LLAMA_MTP_EAGER_CATCHUP=1 restores it for A/B.

Doing that exposed 1e1409a: the QWEN35 MTP builder set t_h_nextn before the inp_out_ids gather, while llama_context::decode copies the first n_outputs rows of it in masked mode. A 4-row MTP batch with one output row therefore returned row 0's hidden state for the draft row, and the second draft pass ran on the wrong input (acceptance 0.86 -> 0.67 on the code prompt). Every MTP decode before this had n_tokens == n_outputs or no outputs, so it never showed. The trunk graph already gathers before publishing; the MTP head now does the same.

Numbers (RTX 4070 12 GB, sm_89, Bonsai-2-27B PTQ1_0)

llama-bench, stock clocks, paired alternating runs:

	prism 9a9394a	#218	#215 (+#220)	this branch (path that dispatches)
tg128	51.0	54.9	58.5	58.5
pp1	47.5	50.5	54.0	54.0
pp2	65.0	92.2	73.8	92.2
pp4	73.0	148.9	93.2	148.9
pp8	93.9	148.3	229.2	229.2
pp512	600	599	1259	1259

Last column: the component measurement of whichever path the hybrid dispatch selects at that column count (SOA_ISUM at 1, PT at 2-4, MMQ from 5). Measured on this exact branch with the +1500 MHz memory clock used for serving: pp1 61.5, pp3 169.5, pp4 212.2 (pp3 145.4 / pp4 167.1 before 847a323).

llama-server, MTP draft (--spec-type draft-mtp --spec-draft-n-max 2), -ctk/-ctv q4_0, one slot, --temp 0, three prompts (code / prose / bash), GDDR6X +1500 MHz:

context	ubatch	mean tok/s	notes
163,840	512	87.6	code 104.4, prose 73.4, bash 86.2; 64.2 with the draft head off
262,144	512	83.7	266 MiB WDDM spill measured via GPU Process Memory counters
262,144	128	86.1	spill gone, 11.5 GB dedicated, full native window resident (before 847a323)
262,144	128	96.6	847a323 + --backend-sampling; code 113.0, bash 96.4, prose 79.9
262,144	128	103.3	+ d8a7180; code 120.9, bash 103.3, prose 86.2; 22.4 ms/step; two fresh processes agree to 0.1; 59.8 with the draft head off

Without e5cd0f3 the 262k + MTP configuration does not fit: nvidia-smi caps at 11.9 GB and WDDM silently pages, 41 tok/s.

Correctness

• test-backend-ops -b CUDA0 -o MUL_MAT: 1283/1283 (45 PTQ1_0 shapes vs CPU); -o MUL_MAT_ID: all PTQ1_0 shapes; -o GATED_DELTA_NET: 39/39.
• test-backend-ops -b CUDA0 -o FLASH_ATTN_EXT: 2984/2984, including new cases added in e5cd0f3 for q4_0/q8_0 K/V at D=128/256, GQA, batch 1..64, KV lengths that are not tile multiples, and permuted layouts.
• With GGML_CUDA_BATCH_INVARIANT=1, 300-token greedy output is byte-identical with the draft head on and off for all three prompts (verify_identity.py from CUDA: fast, batch-invariant small-batch PTQ1_0 mat-vec (1.5x decode on Ampere) #218).
• Perplexity unchanged from cuda: branch-free PTQ1_0 MMQ tile loader and full Ampere tile table (2x prefill) #214 (7.6740 vs 7.6742 stock, float noise).

Scope

CUDA only. HIP keeps the block_q8_1 layout and the old paths (ptq1_0_pt_enabled() is false, ggml_cuda_q8_1_layout_for returns AOS). The in-place K/V flash attention is gated to the MMA kernel on NVIDIA with the head sizes above; other sizes and the vector kernels are untouched.

Repro, recipes and the profiler used to find each step: https://github.com/professorpalmer/bonsai-ada-surgery

[professorpalmer]

Up to 103.3 mean!

[professorpalmer]

Two more commits on top of d8a7180:

• c4922e3 cuda: the Hadamard transform quantizes its own output when every consumer is a PTQ1_0 mat-vec (~390 fewer launches per decode step on Bonsai 2). While soaking it I hit a real hazard worth flagging for anyone fusing producer/consumer here: ggml-alloc hands the transform output the same block as its input on the output-projection path (x is a reshape view of a cont; no in-place reuse through a view parent, so the cont is released at the MUL and the mat-mul takes it two nodes later). The old two-launch path tolerates that; a single fused launch races its own reads and, on a hybrid model, the garbage lands in the recurrent state and poisons the rest of the sequence (surfaced as INT_MAX out of the backend argmax over an all-NaN logits row, about once per 170 requests). The pass now detects the overlap and writes the rows to a pool block held for the graph, and the registry is keyed by tensor identity rather than data pointer. 340 consecutive requests at 262144 ctx with zero errors after the fix.
• 5204d77 sampling: out-of-vocab ids from the backend sampler / draft are logged and rejected instead of reaching vector::at() in the tokenizer.

RTX 4070 12 GB, 262144 ctx, q4_0 KV, draft-mtp n_max=2, backend sampling: 101.6 tok/s mean over 136 requests on the three probe prompts (byte-identical output to the unfused path).

[AlexGabbia]

Short cross-reference for anyone building this integration branch: the one-hunk compile fix posted on #218 (mmvq-ptq1_0.cuh:283-285, GGML_CUDA_RESTRICT on the kernel parameters) applies here unchanged. Without it, d8a7180 fails with 24 × C2912 on MSVC 14.44 (and with GCC 13.3 as reported on #218); with it, the tree builds 411/411 for CMAKE_CUDA_ARCHITECTURES=120 on CUDA 13.0.

Reproduced numbers on an RTX 5070 Ti Laptop (sm_120, 12 GB), paired against the stock fork binary with the same PTQ1_0 file, -ngl 99 -fa 1 -ctk q4_0 -ctv q4_0 -t 8: pp512 519.4 → 1192.3 t/s, tg128 48.0 → 62.9 t/s at 0 depth; pp512 408.9 → 834.6, tg128 34.0 → 42.1 at 32768 depth. With the MTP head at 131072 (q4_0 K/V, --spec-draft-n-max 2): 67.4 → 74.1 t/s on a code prompt, acceptance 0.71.

One thing you may already have covered: #189 (short-query MMA FA with quantized V) is the same shape as the MTP verify batch here (n_max=2 → 3 queries). Worth confirming it is in the integration branch before serving MTP with a quantized cache.

[AlexGabbia]

Follow-up after building the newer head (0974424, i.e. on top of d8a7180 with c4922e3 + 5204d77 + 7e25a91 + 2488ecd + 0974424), same machine, paired runs at a 140 W power limit, same PTQ1_0 and PQ2_0-MTP files:

Plain decode/prefill: flat. llama-bench -ngl 99 -fa 1 -ctk q4_0 -ctv q4_0 -t 8, RTX 5070 Ti Laptop (sm_120):

	d8a7180	0974424
PTQ1_0 pp512 (d0)	1210.2	1198.5
PTQ1_0 tg128 (d0)	63.4	64.2
PTQ1_0 pp512 @ d32768	858.9	846.8
PTQ1_0 tg128 @ d32768	42.8	42.6
PQ2_0 pp512 / tg128 (d0)	1175.3 / 51.8	1183.9 / 51.6

All within ±1%, i.e. run-to-run noise. That is what the bandwidth-bound argument predicts: at batch 1 on this card decode reads ~5.5 GB of weights per token and is saturated on DRAM, so removing ~390 launches per step (the Hadamard-fusion commit) does not show up. The launch-count win presumably needs a card that is launch/latency-bound at batch 1 rather than bandwidth-bound.

MTP: +8%, same policy. llama-server -c 131072, q4_0 K/V, --spec-type draft-mtp --spec-draft-n-max 2, identical prompt, one run per arm: 68.39 → 74.05 t/s decode, with identical draft accounting (149 accepted / 210 generated, mean len 2.42) and identical output length. So the gain is in the speculative loop rather than in the graph as a whole — consistent with the stale-deferred-row and depth-max commits. One run per arm, so treat the +8% as indicative; the direction is credible because the draft accounting is identical.

Net: I would take this head for correctness (the aliasing hazard, the out-of-vocab ids, the stale deferred rows) rather than for speed, and I would measure the fusion commit separately on a latency-bound GPU before claiming a decode win there.

[professorpalmer]

@AlexGabbia a few of those follow-ups are scoring the wrong delta, or attributing commits that do not run in the test you ran.

d8a7180 vs 0974424 on llama-bench looking flat is the expected result, not a verdict on the PR. Those five commits are: fused Hadamard→q8 (correctness hazard + a small launch cut), reject out-of-vocab ids, drop stale deferred draft rows, LIFO pool teardown, and --spec-draft-depth-max. llama-bench never drafts and never crosses a server slot, so four of the five are idle and the fifth is a ~390-launch cut that only shows on a launch-bound card. We already measured that on a desktop 4070 at 23.3 → 23.0 ms/step served, not on tg128. Your 140 W 5070 Ti laptop being DRAM-bound at batch 1 is consistent with a flat bench; it does not mean "take this head for correctness rather than speed."

Speed vs stock on your card is the first comment: pp512 519 → 1192, tg128 48 → 63. That is already in d8a7180. On a desktop 4070 the same stack is 54 → 97 tok/s served (100.7 with the thinking budget off), 632 → 1304 pp2048. Please do not rebase the integration story onto the last-five-commits pair.

The MTP 68.39 → 74.05 (+8%) is not stale-deferred-rows or depth-max. One prompt, one run, no slot reuse, --spec-draft-depth-max default 0: those two commits do not execute. Deferred catch-up is already in d8a7180. Treat a single-run +8% as noise; draft accounting matching does not make the attribution hold.

#189 is not in this branch. It is a separate open PR (MMA FA ncols=32 IMA at Q length 3–4 with quantized V). The in-place q4_0/q8_0 FA path here is a different change (no F16 scratch at prefill). MTP n_max=2 is three queries, so if that IMA still fires on this tree we want a repro — it has not shown up on the 4070 with q8_0 or q4_0 K/V. Not cherry-picked, not claimed.

Recipe, so the 74 tok/s / 128k / q4_0 number does not become the checkpoint: current 12 GB default is 96k / q8_0, draft to 24k (--spec-draft-depth-max 24576). Full trained window is 262,144. 128k/q4_0 with draft always on is the old checkpoint and pages on desktop 12 GB at depth (11.96 GB). Receipts: https://github.com/professorpalmer/bonsai-ada-surgery

[professorpalmer]

Same restrict / C2912 note as on #218, for anyone building this integration tree with CMAKE_CUDA_ARCHITECTURES=120:

The hunk @AlexGabbia posted (GGML_CUDA_RESTRICT off the mul_mat_vec_ptq1_0_pt parameters, onto local aliases) is a real compile fix. It is not in 0974424 yet. It is also not a runtime or speed issue — same kernels, same logits.

We hit that exact C2912 on Windows (CUDA 13, MSVC 14.44) while making the multi-arch zip and did not wait on it: shipped sm_75/86/89 + compute_89 PTX, dropped native 120. 50-series cards JIT the 89 PTX. The 4070 receipts and the live 12 GB recipe do not depend on this hunk.

It should still go in so a clean -DCMAKE_CUDA_ARCHITECTURES=120 tree builds without a drive-by patch. That is the only outstanding code change from this thread.

[professorpalmer]

Serving-recipe note, not a kernel change and not this diff.

The GGUF template defaults reasoning_effort to xhigh (extra "think carefully..." system line). That is the empty-SVG / half-KV-before-first-code-line failure people are hitting. medium is thinking with no extra instruction and is the setting that closed the MBPP/HumanEval gap to the 27B teacher if the output cap is >= 20k; a 10k cap makes medium worse than thinking off. Client max_tokens counts <think>, so a 4k OpenAI default blanks drawings even when the server budget is 20k.

Recorded on the integration binaries, original PTQ1_0, 4 SVG prompts + one short code task, request max_tokens=4096:

arm	empty	SVG	code
think-off	0	4/4	yes
low	4	0/4	yes
medium	1	4/4	no (thought through the cap)
xhigh	4	0/4	yes

Bundle recipe (https://github.com/professorpalmer/bonsai-ada-surgery/blob/main/docs/QUALITY.md): chat medium + --reasoning-budget 20480, agents enable_thinking=false. No change to the CUDA in this PR.

[professorpalmer]

Pushed 229cc80 on this branch:

• GGML_CUDA_RESTRICT off the formal parameters of mul_mat_vec_ptq1_0_pt, aliases inside the body. Same hunk as 2578fdf on CUDA: fast, batch-invariant small-batch PTQ1_0 mat-vec (1.5x decode on Ampere) #218. CMAKE_CUDA_ARCHITECTURES=120 should build without a side patch.
• Ampere (NVIDIA cc 800-889) one-column decode now takes the CUDA: fast, batch-invariant small-batch PTQ1_0 mat-vec (1.5x decode on Ampere) #218 planar kernel. Ada and newer keep the cuda: PTQ1_0 decode on consumer Ampere/Ada: SoA q8 activations, warp-per-row small-K GEMV (+10.5% TG, +14.7% with #220) #215 SoA path. Matches @sudoingX's 3060 table (CUDA: fast, batch-invariant small-batch PTQ1_0 mat-vec (1.5x decode on Ampere) #218 +5.9% tg128 vs cuda: PTQ1_0 decode on consumer Ampere/Ada: SoA q8 activations, warp-per-row small-K GEMV (+10.5% TG, +14.7% with #220) #215 at one column). GGML_CUDA_BATCH_INVARIANT=1 still forces PT at every column count.

4070 numbers are unchanged (still SoA at 1-col). No new zip; this is source on bonsai-combo.

[Copilot]

Copilot review overview

🟡 Changes recommended

Native K/V eligibility and deferred-decode error handling contain unresolved correctness and platform-scope issues.

Get a fresh assessment by requesting another Copilot review.

Review effort: Balanced
Findings: 3 · 3

Open (6)

• Handle failed deferred flush before continuing draft initialization · New
• Gate native quantized K/V dispatch to NVIDIA devices · New
• Check tensor base alignment for quantized loader eligibility · New
• Remove UTF-8 BOM from source file · New
• Correct misleading FWHT fusion function documentation · New
• Remove UTF-8 BOM from new source file · New

What changed in this PR

Integrates CUDA optimizations for full-context Bonsai 2 inference, including PTQ1_0 dispatch, native quantized FlashAttention, recurrent-state fusion, and deferred MTP catch-up.

Changes:

• Adds optimized PTQ1_0 quantization, MMVQ/MMQ, FWHT, and GDN paths.
• Reads q4_0/q8_0 K/V directly in FlashAttention MMA kernels.
• Improves Qwen35 MTP state handling, draft scheduling, and depth limits.

File	Description
tools/​server/​server-context.cpp	Adds speculative depth cutoff.
tests/​test-backend-ops.cpp	Adds FlashAttention and projection cases.
src/​models/​qwen35.cpp	Fixes MTP embeddings and output rows.
ggml/​src/​ggml-cuda/​vecdotq.cuh	Updates PTQ1_0 vector dot logic.
ggml/​src/​ggml-cuda/​template-instances/​generate_cu_files.py	Generates quantized K/V kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q4_0-ncols1_1-ncols2_8.cu	Instantiates q4_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q4_0-ncols1_2-ncols2_4.cu	Instantiates q4_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q4_0-ncols1_2-ncols2_8.cu	Instantiates q4_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q4_0-ncols1_4-ncols2_2.cu	Instantiates q4_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q4_0-ncols1_4-ncols2_4.cu	Instantiates q4_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q4_0-ncols1_4-ncols2_8.cu	Instantiates q4_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q4_0-ncols1_8-ncols2_1.cu	Instantiates q4_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q4_0-ncols1_8-ncols2_2.cu	Instantiates q4_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q4_0-ncols1_8-ncols2_4.cu	Instantiates q4_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q4_0-ncols1_8-ncols2_8.cu	Instantiates q4_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q4_0-ncols1_16-ncols2_1.cu	Instantiates q4_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q4_0-ncols1_16-ncols2_2.cu	Instantiates q4_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q4_0-ncols1_16-ncols2_4.cu	Instantiates q4_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q4_0-ncols1_32-ncols2_1.cu	Instantiates q4_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q4_0-ncols1_32-ncols2_2.cu	Instantiates q4_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q4_0-ncols1_64-ncols2_1.cu	Instantiates q4_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q8_0-ncols1_1-ncols2_8.cu	Instantiates q8_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q8_0-ncols1_2-ncols2_4.cu	Instantiates q8_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q8_0-ncols1_2-ncols2_8.cu	Instantiates q8_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q8_0-ncols1_4-ncols2_2.cu	Instantiates q8_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q8_0-ncols1_4-ncols2_4.cu	Instantiates q8_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q8_0-ncols1_4-ncols2_8.cu	Instantiates q8_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q8_0-ncols1_8-ncols2_1.cu	Instantiates q8_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q8_0-ncols1_8-ncols2_2.cu	Instantiates q8_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q8_0-ncols1_8-ncols2_4.cu	Instantiates q8_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q8_0-ncols1_8-ncols2_8.cu	Instantiates q8_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q8_0-ncols1_16-ncols2_1.cu	Instantiates q8_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q8_0-ncols1_16-ncols2_2.cu	Instantiates q8_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q8_0-ncols1_16-ncols2_4.cu	Instantiates q8_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q8_0-ncols1_32-ncols2_1.cu	Instantiates q8_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q8_0-ncols1_32-ncols2_2.cu	Instantiates q8_0 MMA kernels.
ggml/​src/​ggml-cuda/​template-instances/​fattn-mma-f16-instance-q8_0-ncols1_64-ncols2_1.cu	Instantiates q8_0 MMA kernels.
ggml/​src/​ggml-cuda/​quantize.cuh	Extends quantization interfaces.
ggml/​src/​ggml-cuda/​quantize.cu	Adds PT layouts and fused FWHT quantization.
ggml/​src/​ggml-cuda/​mmvq.cu	Adds hybrid PTQ1_0 dispatch.
ggml/​src/​ggml-cuda/​mmvq-ptq1_0.cuh	Implements planar PTQ1_0 mat-vec.
ggml/​src/​ggml-cuda/​mmvf.cu	Extends small-N mat-vec dispatch.
ggml/​src/​ggml-cuda/​mmq-load-tiles.cuh	Makes PTQ1_0 tile loading uniform.
ggml/​src/​ggml-cuda/​mmq-config-ampere.cuh	Adds wider Ampere tile configurations.
ggml/​src/​ggml-cuda/​ggml-cuda.cu	Adds FWHT and GDN graph fusions.
ggml/​src/​ggml-cuda/​gated_delta_net.cuh	Documents fused gather integration.
ggml/​src/​ggml-cuda/​gated_delta_net.cu	Adds direct state gathering and wider warps.
ggml/​src/​ggml-cuda/​fwht.cuh	Extracts the reusable FWHT butterfly.
ggml/​src/​ggml-cuda/​fwht.cu	Reuses the shared FWHT implementation.
ggml/​src/​ggml-cuda/​fattn.cu	Dispatches native quantized K/V kernels.
ggml/​src/​ggml-cuda/​fattn-mma-f16.cuh	Implements in-place K/V dequantization.
ggml/​src/​ggml-cuda/​fattn-common.cuh	Stabilizes invariant KV splitting.
ggml/​src/​ggml-cuda/​common.cuh	Adds layouts, registries, and feature flags.
common/​speculative.cpp	Defers and merges MTP catch-up work.
common/​sampling.cpp	Diagnoses invalid sampled tokens.
common/​common.h	Adds speculative depth configuration.
common/​arg.cpp	Exposes the depth-limit option.

---

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

[professorpalmer]

Fixed in 3aa89d8. Failed flush_deferred now sets catchup_failed and clears drafting for that sequence in begin() and draft(). The other two call sites already returned false.

[professorpalmer]

Fixed in 3aa89d8. ggml_cuda_fattn_mma_kv_native_supported is now false on HIP/MUSA and on non-NVIDIA cc, so MFMA/WMMA keep the F16 scratch path.

[professorpalmer]

Fixed in 3aa89d8. Eligibility is ggml_cuda_is_aligned(K, 4) && ggml_cuda_is_aligned(V, 4), which checks the base pointer and the higher strides.

[professorpalmer]

Removed the BOM in 3aa89d8.

[professorpalmer]

Removed the leftover GDN preamble in 3aa89d8. The comment now describes only the FWHT-q8 fusion.

[professorpalmer]

Removed the BOM in 3aa89d8.

[professorpalmer]

Serving-recipe note, not a kernel change ? leaving it here because this is the combo / draft-mtp PR.

We grafted ProCreations' on-policy MTP head (the blk.64 tensors only) onto official PTQ1_0. Their combined PQ2_0-MTP GGUF is the wrong pack for Ada and is not what we load.

Same 4070, 96k / q8_0, draft-mtp n-max 2, think off, paired probe:

• teacher Q4_K graft: 93.5 tok/s mean, 66.3% accept (1594/2403)
• on-policy Q8 head: 97.1 tok/s mean, 70.6% accept (1705/2416)

+3.9% decode, +4.3 pp acceptance, +92 MiB. Strip of the graft hashes equal the original PTQ1_0 trunk. Greedy draft-on vs draft-off matches on code and shows the same close-token mmap flip the teacher graft already had, so this is not a new identity regression.

Useful recipe tweak, not a new class of speedup. Default in professorpalmer/bonsai-ada-surgery (make_mtp_procreations.ps1). Nothing for this PR's CUDA diffs.

[bri-prism]

Review of 3aa89d8: [P2] Bring over the corrected PTQ1 shared-memory guard from #218, accounting for this branch's extra padded element.

In ggml/src/ggml-cuda/mmvq-ptq1_0.cuh:459-460 the guard still accepts 2*(K/128)ncols <= 4096. The launcher at line 433 actually requests ncolsrows_per_cta*(K/128+1)sizeof(float)(has_gate?2:1). For the one-column gated path with 64 output rows, K=196608 passes the guard but requests 49184 bytes, exceeding the default 49152-byte block limit. K=262144 also passes and requests 65568 bytes. This affects the enabled planar path, including Ampere one-column dispatch.

Validation: compiled the exact host rows_per_cta helper from this head and checked these launch-byte calculations. This is source/host validation, not a fresh GPU reproduction. #218 already fixed the corresponding unpadded kernel after a reported invalid-argument launch failure. Please share the actual byte calculation between guard and launcher, including the +1 padding here, and add boundary coverage for fused gates.

[professorpalmer]

Thank you for the host-side check. The P2 is correct, and it is on this branch because of the extra pad.

The launcher here requests ncols * rows_per_cta * (K/128 + 1) * 4 bytes, times two with a gate. The +1 is the odd partials stride the epilogue uses for bank-conflict-free pair reads. The old entry test still billed 2 * (K/128) * ncols floats, so a one-column gated CTA at K = 196608 was accepted and then asked for 49184 bytes against the 49152 default. Ampere one-column uses this planar path, so it would have been in range on a wide-K fused shape. No Bonsai 2 27B projection is anywhere near that K.

7e56d13 brings over #218's shared ptq1_0_pt_smem_bytes() and keeps the +1 in that helper, so the guard and the launch count the same bytes and anything over smpb falls back to the generic kernel. Boundary cases in test-backend-ops are shifted by one block versus #218 because of the pad: MUL_MAT at K = 393088 (49152 bytes, fits) / 393216 (fallback), and gated fusion at K = 196480 (fits) / 196608 (fallback).

I have not rebuilt the four-GPU campaign on this head. Happy to if you want it queued. Authorship of the #218 helper is preserved in the comment; this commit is only the pad-aware wiring.

[bri-prism]

Follow-up on 7e56d13: the padded shared-memory guard fix addresses the earlier finding in source; guard and launcher now share the same byte calculation. No fresh CUDA run claimed.

[P2] Reserve capacity for deferred catch-up before appending the first draft row (common/speculative.cpp:2482-2490).

The MTP batch is allocated with llama_n_batch(ctx_dft) entries. process() may defer that many rows when n_tokens <= params.n_max + 1, but draft() appends every deferred row and then one extra anchor without a capacity check. For a single-head, separate-memory MTP context with batch/ubatch 32, n_max=31 and a 32-token final prompt batch, the stash can contain 32 valid rows; the first draft call attempts a 33rd entry and aborts with "llama_batch size exceeded". No rejected-token trimming protects the initial prefill case. Multiple drafting sequences also need their combined catch-up/anchor count budgeted.

Validation: replayed this exact append boundary using the real llama_batch_init/common_batch_add helpers on ARM: 31 catch-up rows plus the anchor succeeds in a capacity-32 batch; 32 plus the anchor aborts (SIGABRT). This is a focused host reproduction plus source-path tracing, not a full model/server reproduction. The previous eager path decoded catch-up separately and did not combine these rows.

Please flush deferred rows before constructing the first draft batch whenever total catch-up plus active draft rows would exceed llama_n_batch(ctx_dft), or split them into bounded decodes. Growing only the host allocation is insufficient because llama_decode also enforces the context batch limit. Add a full-prefill boundary test and a multi-sequence combined-capacity test.

[professorpalmer]

Thank you. This one is right — process() can stash n_batch rows when n_tokens == n_max+1, and draft() then added the first-draft anchor with no remaining slot.

Pushed a59ba13. If catch-up plus one anchor per drafting sequence would exceed llama_n_batch(ctx_dft), we flush the stash first (same decode limit llama_decode enforces), then build the first draft batch. Growing only the host allocation would still abort in decode.

tests/test-mtp-catchup-batch.cpp covers the 32+1 full-prefill boundary and the 16+16+2 multi-seq case. Host-only; no fresh CUDA or server run on this one.

[professorpalmer]

Same MUSA hole as the new finding on #215: once the SoA path was compiled out, the generic PTQ1_0 vec-dot returned zero. 8b3d194 routes MUSA through the HIP scalar implementation. Source fix only; no MUSA hardware here.

[bri-prism]

Reviewed at 8b3d194. I came to this specifically to settle the activation-layout question I deferred on #215 and #218, and it settles it well. The flash-attention work is also in better shape than I expected for a change of that size. Two unreported defects in the speculative path, plus some housekeeping, are what I would want fixed before this lands.

The layout decision is the right one and is built correctly

This is what I asked for on both #215 and #218, which was that the layout be decided here rather than by merge order. ggml_cuda_q8_1_layout_host() is one host-side decision that both the quantizer and the kernel switch consume, the enum documents all three layouts, and the Ampere versus Ada split is measured rather than asserted. Forcing the one-column case onto the planar layout under GGML_CUDA_BATCH_INVARIANT is the right call, since that is the whole point of the flag.

Two notes on it.

The layout is recomputed at three call sites rather than computed once and threaded through, and two of them derive the column count by different expressions (ncols_dst at mmvq.cu:1211 against (int) (ids ? ne2 : ne11) at :1656). The comments at both sites acknowledge that they must match, so the coupling is known, but a future edit to one argument expression breaks agreement silently and nothing would catch it. Computing the layout once in ggml_cuda_mul_mat_vec_q and passing it down would make that structural instead of documented.

ggml_cuda_q8_1_layout_for is marked __host__ __device__ and its comment calls it "the single decision both sides must agree on", but it is not the function that makes the decision: it does not consult cc, so on Ampere it returns SOA_ISUM where the host chooses PT. Nothing calls it from device code today, so this is a trap rather than a bug, but the comment points a future reader at the wrong function. Worth renaming or re-wording.

The FWHT-quantize fusion does the consumer validation properly

Worth saying explicitly because I blocked #187 for the opposite: this matcher walks every use of the transform output, including through reshape views, compares uses_found against ggml_node_get_use_count, and checks input/output overlap. It also requires GGML_OP_MUL_MAT rather than MUL_MAT_ID, which is what makes the hardcoded has_ids = false at ggml-cuda.cu:2895 correct by construction rather than by luck. I went looking for a mismatch there and did not find one.

Flash attention on quantized K/V holds up

I had this audited closely because it is the highest-risk change here, and the things I expected to be wrong were not.

The allocator and the kernel use literally the same predicate function rather than two copies, so get_alloc_size cannot reserve nothing while the kernel takes the conversion path. nstages is forced to 0 for quantized types through matching host and device overloads, so the quantized variants land on the pre-existing non-cp.async path rather than a new one, and every multi-stage block is if constexpr-gated out. The byte-stride change is confined to two lines, and every remaining sizeof(half2) divide sits inside an if constexpr (type_KV == GGML_TYPE_F16) or an nstages > 1 branch. The q4_0 nibble ordering matches the canonical dequant with no swap, q8_0 sign-extends correctly, and the "bit-identical to a to_fp16 pass" claim holds because the product rounds exactly once. The 64 instantiations match the reachable ncols1/ncols2 set exactly.

Three smaller things from that pass:

e5cd0f3 also adds a batch-invariant early return in ggml_cuda_get_best_fattn_kernel (fattn.cu:481-483). That has nothing to do with native quantized K/V, is not in the commit message, and changes kernel selection for batch-invariant builds. It should be its own commit.

HIP and MUSA glob template-instances/fattn-mma*.cu (ggml-hip/CMakeLists.txt:66, ggml-musa/CMakeLists.txt:35), so both build all 32 new instance files, 64 instantiations of a heavy template, for a path whose gate hard-returns false on those backends. Pure compile time and binary size.

The new tests cover 48 cases, but logit_softcap != 0 with a quantized cache is instantiated and never exercised, and max_bias != 0 likewise. At D=128 only two shapes are tested against four plus both permutations at D=256.

Two defects in the speculative deferral path, neither mentioned

Both need a llama_decode failure to reach, so they are low-frequency, but both are real. I verified each in the source rather than taking them on report.

catchup_failed is latched for the life of the object. It is declared at common/speculative.cpp:2084, set at :2112, read at :2447 and :2474, and never assigned false anywhere in the file. begin() resets d.pending and d.n_valid immediately beside it but not this flag, and deferred.assign(n_seq, {}) runs once in the constructor. So a single transient decode failure disables drafting on that slot permanently, for every subsequent task, with no further log. It should be cleared in begin() alongside the other per-task state.

draft() drops the stash before a decode that can fail. d.pending = false is set at :2518 while the batch is being built, and the failure path at :2552 logs and breaks without re-stashing and without setting catchup_failed. Those catch-up rows then exist in neither the stash nor ctx_dft, and the only downstream signal is the soft pos_max < N-1 warning in begin(). Clearing pending after a successful decode, or restoring it on failure, would close it.

Related: the commit described as flushing MTP catch-up "when it plus anchors would overflow n_batch" is fixing a heap buffer overflow that the catch-up batching commit introduced, since batch is sized at exactly llama_n_batch(ctx_dft) while the stash could push one row past it. That is worth naming as an overflow fix rather than a flush condition, because it changes how a reader weighs it.

The 1e1409a MTP fix is genuine and I confirmed the mechanism: the consumer at llama-context.cpp:2197 copies with a literal 0 source offset for n_outputs rows, so before the fix any masked MTP batch with n_tokens > n_outputs read the first catch-up row's hidden state instead of the draft row's. Silent draft corruption, no crash. It is a hard prerequisite for the catch-up batching commit, which is what would have started producing that batch shape, so the ordering matters if anyone cherry-picks from this branch.

Housekeeping before this can land

The branch is CONFLICTING. prism has moved 19 commits since the merge base while this has 25, and the conflicts are in mmq-load-tiles.cuh, src/models/qwen35.cpp and tests/test-backend-ops.cpp, which are exactly the files whose cherry-picks (#214, #217) have since landed independently. It needs a rebase that drops the already-merged commits, and the same will apply to #215 and #218 as they land.

The only CI on this is labeler. For a 60-file CUDA change that includes a flash-attention rewrite, new template instantiations across two backends, and speculative-decoding state machine changes, that is not coverage. This is not this PR's fault and I have raised it separately, but it does mean every correctness claim here rests on the author's own runs plus source review.

Method note: source reading at 8b3d194. I built nothing and ran nothing; the 103 tok/s figure and the per-commit percentages are the author's. Where I say a finding is confirmed, I mean I read the code path myself.

[professorpalmer]

Thank you. Both of those are real, and they are in 1b053d3db.

catchup_failed is cleared at the start of begin(), with the rest of the per-task state. A failed flush no longer disables that slot for every later prompt.

draft() was clearing pending while the batch was still being built, so a failed llama_decode dropped the catch-up rows from both the stash and ctx_dft. pending now stays set until that first decode succeeds. On failure the slot is marked catchup_failed so it is not drafted again until the next begin(). flush_deferred waits for a successful decode the same way.

Rebased onto current prism (bdc23b56). Dropped the two commits that had already landed on their own: the qwen35 Hadamard inverse (#205) and the branch-free MMQ tile loader (#214). No source conflict markers left.

Also reworded ggml_cuda_q8_1_layout_for. It is the column-count helper. ggml_cuda_q8_1_layout_host is the decision both sides call, and it is the one that consults the compute capability. I did not thread that result from ggml_cuda_mul_mat_vec_q yet; the three call sites still recompute it.

Built this head on Windows, sm_89, VS 2022. The first clean build failed: ggml_cuda_q8_1_layout_host was above the declaration of ggml_cuda_info(), so nvcc rejected it in every CUDA file. 18f113d8d moves the helper below that declaration. The CUDA library and llama-common, including these speculative changes, then built clean, and tests/test-mtp-catchup-batch printed ok. 31 catch-up rows plus one anchor fit in a batch of 32, 32 plus the anchor does not, and the two-sequence cases match. That test is the n_batch bound. The two new branches run only when llama_decode returns an error, so this was not a model reproduction of a failed decode.

The flash-attention notes (the batch-invariant early return belonging in its own commit, HIP/MUSA still compiling the new instances, and the untested softcap/bias shapes) are still as they were. Happy to take those next if you want them before this lands.

[AlexGabbia]

Independent sm_120 (Blackwell) build and test check at 18f113d

Windows 11, MSVC 14.44.35207 (VS 2022 17.14), CUDA 13.0.88, CMAKE_CUDA_ARCHITECTURES=120 (promoted to 120a), Ninja, -DGGML_CUDA=ON -DGGML_CUDA_FA_ALL_QUANTS=OFF -DGGML_AVX_VNNI=ON -DLLAMA_CURL=OFF -DLLAMA_BUILD_TESTS=ON -DLLAMA_BUILD_SERVER=ON, clean tree at 18f113d, no side patches:

[513/513] Linking CXX executable bin\llama-server.exe     (0 errors)
test-mtp-catchup-batch.exe                                (ok, exit 0)

tests/CMakeLists.txt:304 registers the test, so it is a real target and not just a source file. The binary runs on the target hardware: llama-server -ngl 99 -c 4096 -fa 1 with Ternary-Bonsai-2-27B-Abliterated-PQ2_0-MTP.gguf, /health 200 and a chat completion returns (system_fingerprint: b341-18f113d, RTX 5070 Ti Laptop, cc 12.0). Every warning in this build is pre-existing and outside the diff: 10039 nvcc #177-D lines from mmq.cuh and mmq-vec-dot.cuh, and 12 MSVC C4244/C4834 lines in src/; nothing originates in fattn-mma-f16.cuh, fattn.cu, speculative.cpp or qwen35.cpp.

Both speculative defects are fixed at this head, read in the source rather than taken from the commit message:

• catchup_failed is cleared at the start of begin() (common/speculative.cpp:2229), with the rest of the per-task state;
• draft() clears pending only after a successful decode (:2573); the failure path sets catchup_failed (:2564) and leaves pending set, so the stash survives and the slot is not drafted again until the next begin().

Stated explicitly because it is the only new test here: test-mtp-catchup-batch asserts the host bound only (catchup + anchors <= n_batch, plus the two degenerate cases). Both defects need a failed llama_decode to reach, so the check above is a source reading, not a reproduction.

Still open at this head, for whoever takes them: the batch-invariant early return is still part of the native-K/V commit (fattn.cu:502); HIP and MUSA still glob template-instances/fattn-mma*.cu (ggml-hip/CMakeLists.txt:66, ggml-musa/CMakeLists.txt:35) while ggml_cuda_fattn_mma_kv_native_supported returns false on both (fattn-mma-f16.cuh:2222); and logit_softcap != 0 / max_bias != 0 with a quantized cache are still instantiated but never exercised.

[professorpalmer]

Rebased onto current prism (842b188, includes #211 / #216 / #257 / #263). Same HIP/MUSA split as #215: HIP keeps the #211 amdgcn_perm body, MUSA stays on the portable scalar loop, NVIDIA still uses _multi. Head is 75b0e43.

The two speculative defects you asked for are still in (catchup_failed cleared in begin(), pending cleared only after a successful first decode). Host bound check reprinted ok: 31+1 fits in 32, 32+1 does not, two-sequence cases match.

Remaining notes from the 8b3d194 / 18f113d reviews, now in 75b0e43:

• The q8 layout is computed once in ggml_cuda_mul_mat_vec_q (and the split-GPU entry) and passed into mul_mat_vec_q_switch_type / switch_ncols_dst. The kernel switch no longer recomputes from ncols_dst.
• HIP and MUSA no longer glob fattn-mma-f16-instance-q4_0-* / q8_0-*. ggml_cuda_fattn_mma_kv_native_supported is still false there; those 32 files were compile-only.
• Isolated ALiBi (max_bias=8) and Gemma softcap (logit_softcap=10) cases on the native q4_0/q8_0 path at D=128, plus ALiBi at D=256, including the permuted D=128 pair.

The batch-invariant vector-kernel early return is still in the native-K/V commit. I did not rewrite that history on the rebase. The comment on it already says it is a batch-invariance policy, not part of the quantized loader.

I did not rebuild CUDA on this box for this push (live 4070 serve stayed up). Happy to queue test-backend-ops -o FLASH_ATTN_EXT,MUL_MAT,MUL_MAT_ID -p q4_0,q8_0,ptq1_0 if you want it before merge.

[professorpalmer]

Picked up @LamplighterPaul's finding from #218.

mul_mat_vec_ptq1_0_pt is launched through ggml_cuda_kernel_launch (PDL on Hopper/Blackwell) and never called ggml_cuda_pdl_sync(), so it can read vy before the q8_1 quantize kernel has written it. Ampere and Ada never take that path; this 4070 box cannot reproduce it. llama-bench / test-backend-ops also miss it, as reported.

Same one-line wait as mul_mat_vec_q, at the top of the kernel, after the restrict aliases. Head is now on bonsai-combo. No speed claim; the 5080 isolation (garbage vs GGML_CUDA_PDL=0 vs this wait) is theirs.

#215 does not carry this kernel. #218 is the source PR; the patch itself is sudoingX#1.

[professorpalmer]

Follow-up from the 5080 notes (thanks @LamplighterPaul).

The 104.8 tg128 / 131K MTP numbers are a win on the full combo, not a miss. Two small adjustments landed on this branch:

1. GGML_CUDA_BATCH_INVARIANT=1 now takes the pre-5300cd1 warp-reduce epilogue on the planar PTQ1_0 mat-vec (smem still uses the bpr+1 stride). Default stays the four-accumulator path. That is the flag our serve already sets for MTP; it should make greedy MTP-on match MTP-off again. The "coastal waters" / "coastal areas" flip was that epilogue, not PDL.
2. ggml_cuda_pdl_sync() in GDN now runs before the s_ids[sequence] load. Safe today because inp_s_copy is a host upload; this is just the wait-before-read rule.

Ada never takes PDL, so I cannot replay the 5080 identity check here. A GGML_CUDA_BATCH_INVARIANT=1 rematch on the 5080 would confirm the restore.

[professorpalmer]

Follow-up measured on the same 4070 (same clocks as the table above), now up as a draft stacked on this branch: #285. Three things in it bear on this PR:

1. --spec-draft-depth-max (319bb14) is superseded. The 24k cutoff was right with deep verify batches on the vector FA kernel. With quantized-KV GQA decode routed to the in-place MMA kernel from e5cd0f3 (one K/V read per KV head instead of per Q head; Bonsai 2 is 24 over 4) and a sliding draft window that keeps the MTP cache small, drafting pays at every depth: 32k 54.5 -> 103.6 tok/s, 64k 48.0 -> 90.1, acceptance unchanged. The MMA routing alone is +11% at 4k and +44% at 16k.
2. The byte-identity claim needs a depth qualifier. Past ~32k, greedy draft-on is not always byte-identical to draft-off with GGML_CUDA_BATCH_INVARIANT=1, and that includes the kernels in this PR as they are: at 40k a code continuation matched, a prose one diverged (both coherent). The FA KV split follows the padded KV length (stream-k partition / parallel-block count), which a verify batch can move one 256-tile ahead of a later single decode; on the vector path the split also depended on per-instance occupancy (fixed in cuda+server: q8_0 K/V at the full 262k window on 12 GB (tiered VMM KV cache), MTP drafting at every depth #285 under the flag). The 300-token check above is unaffected.
3. q8_0 at the full 262k on 12 GB, which this branch could only do at q4_0: a tiered CUDA VMM KV cache (VRAM head, pinned-host tail in the same range, bit-identical output) plus copy-engine staging of the tail. 4k 86.9 / 64k 90.3 / 131k 34.3 / 258k 14.1 tok/s decode.

#215 and #220 are not touched by it (the op-level profile has the PTQ1_0 mat-vecs at ~82% of DRAM bandwidth for single-token decode and the 4-column MMQ crossover here holds up).

[bri-prism]

Independent RTX 4000 Ada (sm_89, CUDA 12.8) check at the current head versus this PR's base. Both were clean CUDA builds on the same card. llama-bench used the same public PTQ1_0 GGUF, -ngl 99 -fa 1 -ctk q4_0 -ctv q4_0, one process at a time. Runs were ordered base/PR/PR/base; each invocation had 3 samples. Values below are medians of the last 2 samples from each invocation (4 steady samples per arm).

Shape	Base tok/s	PR tok/s
pp1	42.29	46.22
pp2	60.33	86.43
pp4	68.31	160.56
pp8	205.39	220.03
pp512	1047.34	1049.32
tg128, depth 0	41.72	46.17
tg128, depth 32768	30.77	33.57

CUDA backend tests: MUL_MAT 1501/1501, GATED_DELTA_NET 39/39, FLASH_ATTN_EXT 2994/2994; test-mtp-catchup-batch returned ok. MUL_MAT_ID was 1034/1035: PTQ1_0 at m=70,n=1,k=2048,n_mats=4,n_used=2 had error 0.02850 versus the 0.00050 limit. On 20 focused repeats of that case, the base failed 3 times and the PR failed 4 times. This is an intermittent failure shared with the base, so the correctness suite is not a clean pass; these samples do not establish a change in failure rate.

Both builds initialized -c 262144 and generated a token; that checks context allocation, not a 262K-token prefill. With GGML_CUDA_BATCH_INVARIANT=1, draft-on/off greedy text matched for one 47,859-token prompt and 48 generated tokens (also matched in a short run). This one long-prefix match does not resolve the >32K prompt-dependent divergence noted in #285.

[professorpalmer]

Thank you for the RTX 4000 Ada run and for landing it.

• MUL_MAT_ID PTQ1_0 m=70,n=1,k=2048,n_mats=4,n_used=2: agreed it is pre-existing (3/20 on the base, 4/20 here), not introduced by this PR. An intermittent error of that size on one shape reads like a partial-tile or uninitialised-read problem rather than rounding, but I have not reproduced it yet. I will run the same case in a loop on the 4070 once the card is free and report back with the path it takes; I am not attributing it to anything until then.
• 47,859-token draft-on/off match: consistent with what cuda+server: q8_0 K/V at the full 262k window on 12 GB (tiered VMM KV cache), MTP drafting at every depth #285 describes. The divergence is prompt-dependent (it needs a verify batch to move the padded KV length one 256-tile ahead of a later single decode), so one matching long prompt neither confirms nor rules it out. cuda+server: q8_0 K/V at the full 262k window on 12 GB (tiered VMM KV cache), MTP drafting at every depth #285's description now cites your check.

Follow-up housekeeping now that this is in prism as 5244cea:

• cuda: PTQ1_0 decode on consumer Ampere/Ada: SoA q8 activations, warp-per-row small-K GEMV (+10.5% TG, +14.7% with #220) #215 and cuda: fold the recurrent-state gather into the GATED_DELTA_NET kernel, per-context registry (+3.8% TG on Ada) #220 are fully contained in this squash (the cuda: PTQ1_0 decode on consumer Ampere/Ada: SoA q8 activations, warp-per-row small-K GEMV (+10.5% TG, +14.7% with #220) #215 layout as the Ada one-column path, the cuda: fold the recurrent-state gather into the GATED_DELTA_NET kernel, per-context registry (+3.8% TG on Ada) #220 gather skip with the per-context registry), so I am closing both as landed here.
• cuda+server: q8_0 K/V at the full 262k window on 12 GB (tiered VMM KV cache), MTP drafting at every depth #285 is rebased onto 445fa82 and is now only its own 6 commits.

[professorpalmer]

@bri-prism following up on the MUL_MAT_ID failure, as promised: reproduced and fixed in #295.

• Reproduced on the 4070: your exact case failed 67 of 300 runs (error up to 0.107), on current prism. Only n=1 fails; n=2..8 at the same shape take the dedicated MoE kernel and were clean.
• Cause: with ids and one token the expert slots are channels and stride_col_dst is nrows * n_used, but the store guard in mul_mat_vec_q compares the row against stride_col_dst. With the small-K geometry (several rows per block) and nrows not a multiple of the block, the last block of one slot stores its out-of-range rows into the next slot's first rows. Two blocks write the same addresses, so the result depends on which lands last. Not introduced by this PR; the same guard is in ggml-org master.
• It is not PTQ1_0-specific. New cases with m = 67 and K = 2 blocks fail on q4_0, q5_0, q5_1, iq1_s, iq1_m, iq2_xxs, iq2_xs, iq2_s, iq4_xs and nvfp4 before the fix.
• After the fix: your case 0/300, the new cases 0 over 30 sweeps, MUL_MAT_ID 1085/1085 and MUL_MAT / FLASH_ATTN_EXT / GATED_DELTA_NET / GET_ROWS all pass, default and with GGML_CUDA_BATCH_INVARIANT=1.