qwen4exp: upstream catch-up, sparse-fa, and NextN/MTP support

[giveen]

@TheTom okay i lied more llama.cpp work.

qwen4exp: upstream catch-up, sparse-fa, and NextN/MTP support

Branch: qwen4-catchup, 11 commits on top of feature/turboquant-kv-cache
(merge-base fb2cc35ab)

Scope

Five pieces of work, landed in this order:

1. Upstream catch-up (1 commit) — every qwen4exp fix upstream shipped
   in the ~10 days before this branch started, that this fork's earlier
   manual port had missed.
2. sparse-fa for DSV4/GLM (3 commits) — cherry-picked from upstream,
   reconciled against TurboQuant's own turbo2/3/4 kernels.
3. Initial NextN/MTP port (1 commit) — brings --spec-type draft-mtp
   to qwen4exp for the first time, porting upstream PR 27836 (open,
   unmerged) plus a standalone-checkpoint loading fix the PR didn't have.
4. NextN/MTP bug fixes (5 commits) — the initial port loaded but
   speculative decoding was completely non-functional (crash, then 0%
   acceptance); this is four independent root causes found and fixed
   against a real model, detailed below.
5. CUDA graph keying fix (1 commit) — cherry-picked from a different
   open upstream PR (models: Qwen3.8-Flash-Next MTP ggml-org/llama.cpp#28243), a generic (not qwen4exp-specific) fix to
   how captured CUDA graphs get cached, relevant to any speculative
   decoding workload on this fork.

git log --oneline feature/turboquant-kv-cache..qwen4-catchup
9f8cb2d3d CUDA: key captured graphs by shape, not first-node address
e5a17b5b9 fix: qwen4exp trunk double-gathered output rows, crashing MTP targets
b7f556e80 fix: export qwen4exp MTP's t_h_nextn before the inp_out_ids gather
97686af29 fix: qwen4exp MTP head should reuse the trunk's shared hc_head mixer
3888e18f0 fix: don't reuse the Gemma4-assistant same-position draft path for real NextN heads
7533986a2 fix: cast F16 norm weights to F32 before CUDA broadcast-mul
209a77e55 qwen4exp: add NextN/MTP draft head (--spec-type draft-mtp)
3ddea31f7 CUDA + ggml: add sparse-fa for DSV4/GLM
f3b503950 test: make the FA V-is-view-of-K case a test case parameter
7fd6a7cc0 CUDA: XOR swizzle flash attn K,V smem fp16 tiles
925792c69 qwen4exp: catch up to upstream through 36b101543

1. Upstream catch-up (925792c69)

Ports every upstream qwen4exp (Qwen3.8-Flash-Next) commit from the
prior ~10 days that this fork's earlier manual port had not received:

• reduce graph splits by hoisting the PLE embedding gather out of the
  per-layer loop (model: qwen4exp: reduce number of graph splits ggml-org/llama.cpp#27880)
• sum indexer heads via strided adds instead of transpose+sum_rows
  (qwen4exp: sum the indexer heads by slices ggml-org/llama.cpp#28023)
• support recurrent state rollback for MTP speculative decoding
  (qwen4exp: support recurrent state rollback ggml-org/llama.cpp#28123)
• rewrite QSA sparse-attention block/bias selection: fixes
  NaN-producing bias rows for short sequences, fixes cross-sequence
  block pooling in a unified KV cache, adds mrope duplicate-position
  ranking, fixes a CUDA rms_norm gridDim.y overflow (qwen4exp: follow up fixes ggml-org/llama.cpp#27941)
• indexer cache seq_cp staleness fix, ext.x/ext.y state-restore fix,
  PLE-must-be-linear-attention validation, correct -sm tensor
  disablement (qwen4exp: follow up fixes ggml-org/llama.cpp#27941)

2. sparse-fa for DSV4/GLM (7fd6a7cc0, f3b503950, 3ddea31f7)

Cherry-picked from upstream 8e93a9773, on top of two prerequisites:

• 7fd6a7cc0 — XOR swizzle layout for flash-attn K/V smem fp16 tiles
  (fattn-swizzle.cuh), needed by the sparse gather-load tiles.
  Reconciled against TurboQuant's own turbo2/3/4 SRAM tile loaders in
  fattn-mma-f16.cuh, left untouched.
• f3b503950 — makes the FA V-is-view-of-K test case a named
  constructor parameter instead of a hardcoded bool, and wires
  kv_view through to K/V tensor creation, matching the new sparse-fa
  test cases' needs.
• 3ddea31f7 — the feature itself: ggml_flash_attn_ext_set_n_kv_max()
  lets the CUDA flash-attention kernel treat the mask's finite entries
  as a sparse K/V set and skip the rest, instead of computing dense
  attention and masking it out. DeepSeek-V4's own top-k
  sparse-attention path is wired to it. Reconciled against
  TurboQuant's dedicated fattn-mma-turbo.cuh kernel, which the
  upstream diff had no knowledge of — two real positional-argument
  bugs were caught and fixed here (a use_sparse bool landing in the
  wrong template/call-site slot in two places) before they could ship.
  Confirmed neutral/no-regression via full test suite + benchmark, as
  expected since qwen4exp's attention shape doesn't hit the new
  kernel's gate.

3. Initial NextN/MTP port (209a77e55)

Ports upstream PR ggml-org#27836 (open, unmerged) on top of
this fork's qwen4exp support: the MTP head folds the next token's
embedding into the trunk's wide hyper-connection residual, runs one
trunk-shaped block (dense attention + MoE) over it, and collapses the
result with its own hyper-connection mixer before reusing the trunk's
LM head. Also updates the converter and gguf-py tensor mappings so
conversion/qwen4exp.py --mtp can export the head.

Also adds mtp_only/trunk_flags handling to load_arch_tensors,
which the ported PR didn't include: without it, loading a standalone
MTP-only checkpoint (all trunk tensors absent) fails outright instead
of tolerating their absence, unlike the equivalent qwen35.cpp path.
Detects MTP-only via the absence of blk.0.hc_attn_norm.weight, a
tensor every trunk layer carries.

At this point the port loaded but speculative decoding did not
work — first a hard crash, then (once crash-free) 0% draft
acceptance. The next four commits are the root-cause fixes for that,
found and verified against real models rather than assumed correct.

4. NextN/MTP bug fixes (5 commits)

--spec-type draft-mtp for qwen4exp went from "crashes on load" to
"loads but gives 0% draft acceptance" to "works, with real measured
acceptance in line with the community-validated implementation." Five
commits, four independent root causes, each confirmed against a real
model rather than assumed fixed.

Files changed (this section only)

common/speculative.cpp  |  6 ++++-
src/llama-context.cpp   |  4 ++++
src/llama-ext.h         |  7 ++++++
src/llama-graph.cpp     |  9 ++++++++
src/models/qwen4exp.cpp | 61 ++++++++++++++++++++++++++++++-------------------
5 files changed, 63 insertions(+), 24 deletions(-)

Commits

1. 7533986a2 — cast F16 norm weights to F32 before CUDA broadcast-mul

The original bug report: llama-server aborted on load with
GGML_ASSERT(nb10 % sizeof(src1_t) == 0) in ggml_cuda_op_bin_bcast
(binbcast.cu). Root cause: that CUDA kernel has a dispatch path for
F16-activation × F32-or-F16-weight, but none for F32-activation ×
F16-weight. Real MTP checkpoints store some norm weights
(nextn.enorm, etc.) as F16 — unconventional but valid — while trunk
norm weights happen to be F32 in every checkpoint tested so far, which
is why this never fired outside the MTP path. Fixed in the three call
sites that mix a norm weight into an activation (build_norm,
build_hc_mix, build_ple): cast the weight to F32 first when the
combination would otherwise hit the missing dispatch branch. No-op
when weights are already F32.

2. 3888e18f0 — don't reuse the Gemma4-assistant same-position draft path

The speculative driver's is_mem_shared flag (true whenever no
separate -md draft model is given) was being used to select
Gemma4-assistant's same-llama_pos-for-every-draft-token behavior.
That behavior is specific to Gemma4-assistant's early-exit
self-speculation, not a general property of "shares KV memory with
the target." Added llama_model_uses_shared_position_draft(), gated
on arch == LLM_ARCH_GEMMA4_ASSISTANT specifically, and used it
instead of the blanket is_mem_shared check for that one branch. Real
fix, though it turned out not to be the cause of the 0%-acceptance
issue investigated afterward — self-contained MTP (no -md) on this
architecture happened to resolve is_mem_shared=false anyway.

3. 97686af29 — MTP head should reuse the trunk's shared hc_head mixer

The ported PR (27836) gave the MTP block its own private per-layer
final mixer (layer.nextn.hc_head_norm/down/up,
blk.N.nextn.hc_head_* on disk). The actual, community-validated
implementation (PR 27739, reconciled by LaurentZuijdwijk — what real
MTP-head exports like dzannotti/Qwen3.8-Flash-Next-MTP-GGUF are
built against) instead trains one hc mixer shared between the
trunk's own final layer and the MTP head, exported as top-level
output_hc_norm/down/up and loaded as model.hc_head_*. Confirmed
directly against a real MTP head export: it has no
blk.N.nextn.hc_head_* tensors at all. Switched graph_mtp to use
model.hc_head_*.

4. b7f556e80 + e5a17b5b9 — t_h_nextn export timing and a trunk double-gather

Two related bugs in the hidden-state handoff:

• graph_mtp gathered res_hc down to output-only rows before
  exporting t_h_nextn. Any decode with fewer requested-output rows
  than input rows — concretely, the driver's catch-up/prefill decode,
  which requests logits for none of its rows — exported a zero-row
  hidden state instead of the full per-token state the driver's
  shift-by-one handoff needs. Fixed by exporting before the gather,
  matching PR 27739's ordering.

• The trunk graph (graph::graph, not graph_mtp) has two gather
  paths for res_hc/cur — an early one and a deferred one — gated
  by complementary conditions, so exactly one always fires when
  inp_out_ids exists. A third, unconditional gather right before
  the final output norm re-applied inp_out_ids again, indexing an
  already-reduced tensor with values sized for the original token
  count. Silently harmless whenever n_outputs == n_tokens (why
  normal generation never hit it — a redundant identity gather is a
  no-op), but a hard crash (GGML_ASSERT(i01 >= 0 && i01 < ne01) in
  ggml_compute_forward_get_rows) whenever n_outputs < n_tokens on
  a context with embeddings_nextn set — which is unconditionally
  true for any target paired with an MTP draft. This one crashed on
  literally every server startup with a real -md draft-mtp config,
  independent of which draft checkpoint was used or which graph_mtp
  bugs remained. Fixed by removing the redundant gather.

Also made layer.nextn.hc_head_norm/down/up TENSOR_NOT_REQUIRED
(dead weight as of commit 3, kept optional only so PR-27836-style
exports that still carry it continue to load).

5. CUDA graph keying fix (9f8cb2d3d)

While investigating whether anything else upstream would help, found
PR #28243 (open,
unmerged, built on top of ggml-org#27836), which claims "1.3 to 2x faster MTP
support for Qwen3.8-Flash-Next." Checked its diff line-by-line against
what's already fixed here: it does not fix the mixer-sharing or
export-timing bugs above (both still present in its graph_mtp,
unchanged from ggml-org#27836) — its qwen4exp-specific parts are a different
feature (embed_tokens/output sharing between trunk and draft to
save VRAM/disk) and a clearer error message for a different edge case
(loading a draft-only export standalone instead of via -md).

The actual source of its speed claim is generic CUDA infrastructure,
unrelated to qwen4exp: the CUDA graph cache keyed captured graphs by
the raw memory address of their first node. A captured graph
hard-codes its shapes, but speculative decoding constantly alternates
between different batch shapes (draft steps, verify batches, catch-up
decodes) on the same context — when a new shape happens to reuse the
same first-node address as a stale cached graph for a different shape,
capture thrashes instead of ever reaching steady-state replay. Cherry-
picked just this piece (not the qwen4exp-specific parts of ggml-org#28243):
hash the node count and both endpoint tensors' shapes into the key
instead of the raw pointer, with LRU eviction capped at 64 graphs.

This benefits any speculative-decoding workload on this fork, not just
qwen4exp. Verified no regression (output still byte-identical at
temp 0); its own speed contribution is hard to isolate cleanly from
the MoE-cache warmup effect documented below, which dominates in the
same test runs — kept on correctness grounds (the old keying scheme
had a real cache-key collision problem) independent of the measured
delta.

Validation

Models used

Target: unsloth/Qwen3.8-Flash-Next-GGUF, quant UD-Q4_K_XL (4
shards, ~112 GB), an unmodified quantization of the official
Qwen/Qwen3.8-Flash-Next release — no abliteration/uncensoring.

Draft head: dzannotti/Qwen3.8-Flash-Next-MTP-GGUF, file
Qwen3.8-Flash-Next-MTP-Q4_K_M.gguf (2.62 GB), downloaded to
/mnt/storage/models/qwen3.8/MTP/. This is the MTP block that ships
inside the official Qwen/Qwen3.8-Flash-Next checkpoint but that
upstream's own converter (PR 27742) drops during conversion — this
repo exports it standalone (convert_hf_to_gguf.py --mtp) so it can
be attached to any clean target GGUF via -md, without needing to
requantize the ~180B-parameter full model. Chosen specifically because
it's not built from an abliterated/uncensored merge, unlike the
checkpoint in the original bug report (see below).

Both are the models referenced in run.txt in the same folder as this
document, which has the exact launch command used for the benchmarks
below.

Correctness

• Loads without crashing (previously: hard crash on every startup
  with any real -md draft-mtp config, independent of checkpoint —
  see commit e5a17b5b9).
• 100% draft acceptance on a trivial prompt (26/26).
• 77–84% draft acceptance on natural-language generation, matching
  the community-reported 0.74–0.90 range for this exact head.

The original bug-report checkpoint (a third-party "Uncensored"
abliterated merge, since deleted from local storage) still gets 0%
acceptance with all four fixes applied — that confirms it's the
checkpoint's own weight data at fault, not a code issue.

Speed: MTP on vs off, same target, same hardware, same prompt

All runs: UD-Q4_K_XL target, -c 8192, -fa on, temp 0
(deterministic — MTP-on and MTP-off produce byte-identical output).
MTP runs add -md <head> -ngld 999 --spec-type draft-mtp --spec-draft-n-max 3 --spec-draft-p-min 0.75.

config	tok/s	vs baseline	draft accepted
baseline, no speculative decoding	32.31	—	—
+ MTP, cache cold (fresh server restart)	33.65	+4.2%	88/114 (77.2%)
+ MTP, cache warm (repeated/related prompts)	up to 52.98	+63.9%	230/276 (83.3%)

The cold-cache number is the honest "what you get from a single fresh
request" figure. The warm-cache number is what a sustained session
converges to — see the next section for exactly how and why. Either
way, MTP is a net win on this target/hardware once fixed; it was a net
loss before these fixes (checkpoint-dependent 0% acceptance made the
draft pass pure overhead).

Speed scales with MoE-cache warmth; acceptance does not

Every single-shot benchmark in this session used a freshly-restarted
server, which measures the MoE expert cache cold. Once the cache
warms up under sustained load, the same target+draft pair is
substantially faster — this is a real, measured effect, not a code
change, and worth knowing before trusting any single-run number.

Setup: UD-Q4_K_XL target + dzannotti MTP head via -md, -c 8192,
--spec-draft-n-max 3 --spec-draft-p-min 0.75, --moe-cache auto,
temp 0 (deterministic output), GGML_CUDA_MOE_CACHE_STATS=20 for
live hit-rate telemetry.

Cache fill, same server, three consecutive requests:

run	prompt	tok/s	draft accepted	cache hit rate at request end
1 (cold start)	printing-press essay	44.62	230/273 (84.2%)	climbed 4.4% → 25.0% over the run; used=6998/20998 (33% full), evictions=0
2 (same prompt repeated)	printing-press essay	52.98	230/276 (83.3%)	83.4% steady state; used=20998/20998 (100% full), evictions=10442
3 (different topic)	photosynthesis	30.99	89/111 (80.2%)	still on the same warm server as run 2

Run 2 vs run 1: +18.7% tok/s from cache warmth alone, identical
prompt, identical everything else, output byte-identical (temp 0).
Run 3 shows the warmth is partly topic-specific: a fresh topic on an
already-warm server still lands close to the cold-start number, not
the same-topic peak — some experts are apparently "generalist" and
stay hot across any topic (giving a baseline lift over a truly cold
cache), but a large share of run 2's peak came from that specific
topic's expert subset being fully resident from repetition.

Root cause, confirmed by reading moe-cache.cu: the cache fills and
evicts based on live usage with no pre-warming — evictions=0 in run
1 means it hadn't even reached its capacity limit yet, so the
33%-full 25%-hit-rate snapshot at the end of run 1 is mid-fill, not
steady state. Run 2 shows what steady state (used == full capacity,
evictions actively happening) actually looks like: 83.4%, over 3x the
end-of-run-1 snapshot.

Draft acceptance did not track this. It stayed in the 77–84% band
across all three runs regardless of cache warmth (77.2% earlier in
this session on a cold -c 8192 server, 84.2%/83.3%/80.2% here) — the
variation looks prompt-dependent, not cache-state-dependent, which
makes sense: the MoE cache returns bit-identical expert weights
whether served from cache or a fresh read, so it changes latency,
not the draft head's predictions. Speed and acceptance are two
independent effects; only speed benefits from warmup.

Practical implication: don't judge this setup's real throughput
from a single fresh-restart benchmark — expect a meaningfully higher
number in a long-running session, especially on repeated or related
topics, plateauing once the cache fills and reaches its own steady
per-topic hit rate.

Known follow-ups (not part of this PR)

• --spec-chain (chained in-graph multi-token MTP drafting) is not
  supported for qwen4exp (llama_model_supports_mtp_chain() only
  returns true for LLM_ARCH_QWEN35). Porting it means re-deriving
  qwen4exp's hyper-connection block (attention + MoE + hc-combine) in
  the same low-level, per-row, in-graph, direct-KV-write form qwen35's
  chain path uses — a real feature port, not a quick fix. The
  reference implementation's own author also found deeper/chained
  drafting didn't help this specific model (its draft head carries a
  full MoE, so every drafted token is already a real forward pass).

• MoE cache VRAM sizing uses a static reserve (GGML_CUDA_MOE_CACHE_RESERVE_MB,
  default 3072 MiB) rather than measuring actual runtime peak usage.
  Reducing it crashed with a genuine CUDA OOM on the first real decode
  after a speculative-decoding config was added, meaning something in
  the live MTP serving path (plausibly the cache's own async fill
  queue) grows past the load-time worst-case graph reservation.
  Root-caused to a real gap, not chased to a fix.

  AI usage: yes

[gopinath87607]

hi if you have time kindly check this too

Thanks, Mark. That completes the DS4 hardware matrix.

I pushed 06b66fd with the upstream LIGHTNING_INDEXER Meta handler. It requires all four inputs to be mirrored and returns a mirrored result, matching every occurrence you traced while rejecting any unsupported split input rather than silently accepting it.

Local results on the exact head:

full test-llama-archs: pass
test-turbo-quant: pass
test_slot_save_restore: pass
Release Metal and llama-server builds: pass
git diff --check: clean

Your patched-tree results already cover the important real-hardware cases on this exact rule: coherent output on 2 devices, 4 devices, and uneven 3,4,4,1, including fresh prompts, sequence reset, and prefix reuse. No further DS4 retest is needed unless fresh CI finds something.

I am tracking the long-context Qwen4Exp allocation report separately. Its log shows a 5.96 GiB compute-buffer request failing with --batch-size 6700 --ubatch-size 6700; the subsequent segfault after allocation failure is not an acceptable failure mode, but it is separate from the DS4 split correctness work above.

turbo qunat is the only fork thats support proper split and proper buffer size

[giveen]

./build/bin/llama-server \
  -m /mnt/storage/models/qwen3.8/flash/UD-IQ4_XS/Qwen3.8-Flash-Next-UD-IQ4_XS-00001-of-00003.gguf \
  -md /mnt/storage/models/qwen3.8/MTP/Qwen3.8-Flash-Next-MTP-Q4_K_M.gguf \
  -ngld auto \
  -c 32000 -ngl auto --fit on -fa on --jinja -t 8 -Cr 0-7 -tb 8 -Crb 0-7 \
  --cpu-strict 1 -b 512 -ub 512 -ctv q8_0 -ctk q8_0 \
  --top-p 0.95 --top-k 20 --temp 1.0 --parallel 1 \
  --host 0.0.0.0 --port 8099 --moe-cache auto \
  --spec-type draft-mtp --spec-draft-n-max 3 --spec-draft-p-min 0.75 \
  --chat-template-kwargs '{"reasoning_effort":"medium"}' \
  --chat-template-file /mnt/storage/llama-server/qwen3.6-chat-template/chat_template.jinja \
  --reasoning-preserve

Prompt 1:

• Prompt eval speed: 25.23 tokens/s
• Generation speed: 37.99 tokens/s
• Draft acceptance: 90.68%

Follow up Prompt 2 about improving a python script:

• Prompt eval speed: 26.44 tokens/s
• Generation speed: 32.31 tokens/s
• Draft acceptance: 93.00%

[TheTom]

Thanks for the writeup, it made this reviewable. Most of it holds up: the sparse-fa port is byte-for-byte upstream 8e93a9773 apart from the fork's trailing type_K, type_V template params, both use_sparse slot fixes in fattn-mma-turbo.cuh are correct against the new signatures, the op_params[4] slot does not collide with anything, and the four MTP root-cause fixes check out against the code (the double-gather one in particular: gather_now and the deferred gather really are complementary, and the third gather was redundant). On the Mac (M5 Max, Metal), test-llama-archs passes for qwen4exp/qwen35/qwen35moe, and test-backend-ops on MTL0 gives FLASH_ATTN_EXT 11815/11815 (1364 of those are turbo K/V cases), MUL_MAT_ID 893/893, SET_ROWS 159/159, TURBO_WHT 27/27. MUL_MAT has 15 tq3_1s/tq4_1s sentinel failures, identical on the base branch, that is #285, not this PR.

One blocker, one thing to fix with it, and one bookkeeping item.

Blocker: the XOR-swizzle port breaks turbo2/3/4 flash attention on Turing and newer.

7fd6a7cc0 ports upstream e4b9af007 and changes the shared K/V tile layout in fattn-mma-f16.cuh:

constexpr int  stride_tile_K = ggml_cuda_fattn_smem_swizzle::tile_stride(nbatch_K2);   // was nbatch_K2 + 4
constexpr bool swz_K         = ggml_cuda_fattn_smem_swizzle::enabled(nbatch_K2);       // true when nbatch_K2 % 32 == 0

The f16 loader was updated to write through bytes_rc<>(), but the fork's three writers, flash_attn_ext_turbo{2,3,4}_load_tile (lines 545-702 at your head), still store tile_KV[row*stride_tile + c] linearly, while every read of that tile now goes through ggml_cuda_fattn_smem_swizzle::load_ldmatrix<stride_tile_K, swz_K> (lines 920, 946, 1297, 1323). Write layout and read layout disagree whenever swz_K is true.

That is the default case, not a corner: fattn.cu routes turbo K/V at head dim 128 and 256 with Q->ne[1] <= 4 on turing_mma_available(cc) to ggml_cuda_flash_attn_ext_mma_turbo_case, and ggml_cuda_turbo_mma_fused() returns true unless GGML_TURBO_MMA_FUSED=0 (the "DEFAULT OFF" comment above it is stale). Head dim 128 gives nbatch_K2 = 64, 256 gives 128, both bank-aligned, so every fused turbo decode on an NVIDIA GPU hits the mismatch.

Reproduced on a GB10 (sm_121, driver 580.173.02), your head e5a17b5b9, -DCMAKE_CUDA_ARCHITECTURES=121:

./build/bin/test-backend-ops test -b CUDA0 -o FLASH_ATTN_EXT
  6934/7686 tests passed
  Backend CUDA0: FAIL

All 743 FLASH_ATTN_EXT failures are type_K=turbo3 (368) or turbo4 (375), all hsk=128, all nb=1 or nb=3, with ERR around 1.5 against a 5e-4 tolerance. Not a single turbo case passes. Example:

FLASH_ATTN_EXT(hsk=128,hsv=128,nh=4,nr23=[1,1],kv=113,nb=1,mask=1,sinks=1,prec=f32,type_K=turbo3,type_V=turbo3,...): ERR = 1.570429459 FAIL

Same box, same flags, base branch at 80be9a7: 7658/7658 tests passed, zero turbo failures. So this is a regression introduced here, not a pre-existing GB10 issue.

Fix: make the three turbo writers go through bytes_rc<stride_tile>(row, c) like the f16 loader (or take a swz template param and branch), and update fattn-mma-turbo.cuh:47, which still sizes host SMEM as nbatch_fa * max(nbatch_K2 + 4, nbatch_V2 + 4) * sizeof(half2) while the device kernel now uses tile_stride(). Over-allocation only today, but it should come from the same helper as fattn-mma-f16.cuh:2273. Then add turbo3/turbo4 FLASH_ATTN_EXT cases at hsk 128 and 256 with nb <= 4 and nr2 >= 2 to test-backend-ops so the fused turbo path is provably hit; the four new swizzle-stride cases in this PR are all F16 KV, which is how this got through.

Bookkeeping: the description lists 11 commits including 9f8cb2d3d (CUDA graph keying from upstream ggml-org#28243), but the PR has 10 and ends at e5a17b5b9. There is no ggml-cuda.cu hunk in the diff. Either push it or drop it from the writeup.

Smaller items, fine as follow-ups:

• llama-graph.cpp build_norm: the ggml_cast(mw, F32) runs on every graph build for every arch whose norm weight type differs from the activation, and it also lands on CPU/Metal/Vulkan which never had the missing binbcast.cu dispatch. Cheaper: upcast narrow norm weights once at load, or add the F32-act x F16-weight branch to binbcast.cu.
• deepseek4.cpp:784: n_kv_max = min(mask->ne[0], n_swa) + top_k matches upstream, but nothing enforces the "bounds the finite entries per row" contract; flash_attn_mask_to_sparse_indices silently drops entries past n_kv_max. A debug assert would catch an n_swa == 0 model.
• llama-kv-cache.cpp:2248: if (k_rot && k_rot->buffer) silently skips a fork input when unallocated. If that is expected in some path, a comment saying when; otherwise an assert.
• fattn.cu:190: the "DEFAULT OFF" comment on ggml_cuda_turbo_mma_fused() is wrong, pre-existing but you are in the file.

Once the turbo writers are swizzle-aware and a turbo FA case runs green on CUDA, I'm happy to merge this.

[TheTom]

Requesting changes for the turbo2/3/4 swizzle mismatch detailed in the comment above (743 CUDA FLASH_ATTN_EXT failures on GB10, base passes 7658/7658). Everything else looks good.

[TheTom]

Round two on 69499d4. The blocker is fixed properly: all six turbo store sites (zero-fill and main store for turbo2/3/4) go through turbo_store_h2<stride_tile, swz>(), the call sites pass the same swz_K/swz_V and stride_tile_K/stride_tile_V the readers use, and fattn-mma-turbo.cuh now sizes host SMEM with the same tile_stride(nbatch, cc) helper as fattn-mma-f16.cuh. The graph keying commit is byte-identical to upstream ggml-org#28243's common.cuh / ggml-cuda.cu hunks, and the node_properties re-capture check still runs on a key hit. The new hsk=256 turbo3/turbo4 cases really route to the fused turbo MMA path (nb 1 or 3, nr2 4 forces kv 512, use_gqa_opt holds), so they are not vacuous.

Fork preservation: I diffed every file from the merge-base and checked each deleted line that does not exist in upstream. Nothing fork-specific goes away. The 195 such lines are signature rewrites for the sparse and swizzle params, the turbo loader fix itself, the FA test sweep gaining turbo2 and hsk=256 cases, the fused-turbo comment corrected to match the code (still default ON), and the fork's older hand-port of qwen4exp / QSA replaced by upstream's newer version. Local writeup kept on my side.

GB10 (sm_121), head 69499d4, -DCMAKE_CUDA_ARCHITECTURES=121:

test-backend-ops test -b CUDA0 -o FLASH_ATTN_EXT    8774/8774   (2592 turbo K/V cases, 336 of them at hsk=256; round one had 743 failing)
test-llama-archs                                    all OK
full sweep, graphs on                               segfault at ROLL, same point as base (#347), so 9f8cb2d3d does not fix that one

The no-graphs full sweep is running now as the last gate; I'll merge when it and the remaining CI jobs are green.

Follow-ups, none blocking:

• 69499d4 does more than skip V. Setting n_embd_head_{k,v}_mla_impl makes hparams_idx.is_mla() true, which also turns off TURBO_AUTO_ASYMMETRIC's K upgrade for the indexer cache (llama-kv-cache.cpp around line 151). Before this commit, --cache-type-k turbo{2,3,4} on qwen4exp silently upgraded the indexer K to q8_0 because n_head_kv_arr of 1 gave a meaningless GQA ratio; now it stays turbo. Probably the better behavior, and it matches the DSv4 hparams_lid precedent, but it is a change in a fork knob and should be in the commit message. If you have a ppl number with a turbo K cache on this head, that settles it.
• The hsk=256 turbo2 cases only cover VEC, since fattn.cu deliberately keeps turbo2 off the fused path at 256. A comment on the test would stop someone reading that as swizzle coverage. turbo2's swizzled path is covered at 128.
• fattn.cu around line 928 still has the old sentence fragment "Default OFF (env unset / !=1) falls straight" right before the corrected one, and says "turbo4-K==turbo4-V" where the gate handles turbo2/3/4. I'll tidy that on the base after merge rather than churn your branch.
• I checked the new deepseek4.cpp assert against a small context statically: n_top_k is clamped to indexer_score->ne[0] and the raw half to raw_mask->ne[0], so n_kv_max cannot exceed the concat. Fine as is.

[TheTom]

Merged as b5b813d. For the record, the fork-preservation check in detail, since this was a catch-up and the question was whether any fork-only work got dropped.

Method: three-dot diff from the PR merge-base fb2cc35, and for every file the PR touches, every deleted line that does not exist in upstream/master and is not re-added in the same file. 34 files, 195 such lines, each group read for meaning.

file	lines	what it is	fork behavior kept?
ggml-cuda/fattn-mma-f16.cuh	91	turbo2/3/4 tile writers rewritten to the swizzled layout; f16 loader, mask loader and kernel signatures gain the sparse/swizzle params	yes, this is the round-one fix
ggml-cuda/fattn-mma-turbo.cuh	4	host SMEM sizing moved to tile_stride(); kernel template gains use_sparse	yes
ggml-cuda/fattn-common.cuh, fattn-tile.cuh, fattn-vec.cuh	11	launch_fattn gains use_sparse; all call sites updated	yes
ggml-cuda/fattn.cu	3	comment on ggml_cuda_turbo_mma_fused() corrected from "DEFAULT OFF" to "DEFAULT ON"; code unchanged and still ON	yes
tests/test-backend-ops.cpp	17	FLASH_ATTN_EXT sweep keeps the turbo gating and extends it to turbo2 and hsk=256; old perf loop replaced by upstream's larger one	yes, coverage grew
src/llama-kv-cache.cpp	1	if (k_rot) becomes if (k_rot && k_rot->buffer) with the explanation	yes
src/llama-graph.cpp, deepseek4.cpp, qwen35.cpp	11	build_attn_mha gains n_kv_max; call sites updated	yes
src/llama-memory-hybrid-idx.{cpp,h}, src/models/qwen4exp.cpp	55	the fork's earlier hand-port of upstream QSA/qwen4exp replaced by upstream's newer version; these files hold no turbo/TQ logic before or after	yes, this is the catch-up itself
src/llama-model.cpp, conversion/qwen4exp.py	2	QWEN4EXP added to an arch list; unused import dropped	yes

No GGML_TQ_* or TURBO_* knob, no TQ dispatch, no turbo kernel and no fork test was removed. The 37 base commits this branch did not have (#333 through #344) merged cleanly; the only shared file with real overlap was common/speculative.cpp, where #344's DFlash hunks and this PR's MTP hunks live in different structs.

Final gates on a GB10 (sm_121), head 69499d4: FLASH_ATTN_EXT 8774/8774 with 2592 turbo cases; test-llama-archs clean; full sweep with GGML_CUDA_DISABLE_GRAPHS=1 19512/19512; hosted CI green including the HIP VGPR check. The stale fattn.cu comment I mentioned is fixed on base in cde1762.