Repository navigation
qsa : skip foreign sequences when filling scan tails - #165
Open
iTechMedic wants to merge 388 commits into
Open
iTechMedic wants to merge 388 commits into
iTechMedic wants to merge 388 commits into
Conversation
No behaviour change on any path that runs today; this removes footguns the split path left in shared code and clears the leftovers. - flash_attn.comp and flash_attn_cm2.comp still read p.k_num and p.gqa_ratio raw, while the split path packs a partition count into the high half of k_num and a flag into bit 31 of gqa_ratio. Only the coopmat1 gate keeps them from ever seeing those values, and nothing said so. Use the decoded globals init_indices() already computes, so routing the split path at a sibling shader fails loudly instead of writing at wild offsets. - Drop profile_stage and GGML_VK_FA_TOPK_PROFILE. This also removes two continues from the sparse kernel's key loop. - Remove the unreachable mask-stride override branch (nothing sets bits 16..30 without bit 31) and restore the comment explaining why the GQA case writes zero the roundabout way: the compiler must not fold it, or stride alignment detection breaks. - Assert that the mask stride covers the raw KV range. The raw dispatch smuggles that stride through split_kv, which the shader also uses to derive its KV range, so a stride below KV would silently clip it. - Request descriptor sets only after the split scratch is known to fit, so the fallback path does not inherit n_tiles of unused requests. - n_head_log2 is unread with max_bias == 0; pass 0 rather than a number that looks computed. - Make the crossover depend on the path that will run. The coopmat sparse path wins as soon as any key is pruned, but the scalar fallback does not and regresses against dense FA until the pruned fraction is large, so it keeps its original 3x margin. - Record that the coopmat sparse shader requires 64-wide subgroups and eight of them, at the constants that encode it. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Assisted-by: Claude (Opus 5)
The sparse top-k fixture pinned every tensor's fourth dimension to 1, so no case reached the split path's stream indexing and three separate mis-indexings passed the suite. Add an ns parameter and three cases: two at ns=2 (with and without sinks, single tile) and one at ns=3 with 300 query tokens, which also crosses the 256-token tile boundary. The per-token selection is offset by the stream so a dropped stream stride reads another sequence's keys rather than the same ones. All three fail on the unfixed shaders and pass with them fixed. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Assisted-by: Claude (Opus 5)
flash_attn_split_k_reduce was using one value for two different things: ne2 is the split buffer's query count, which the sparse path hands the TILE height, but the destination's stream stride needs the FULL query count. With several tiles and more than one sequence it wrote stream s at s*tile_size instead of s*N, so the previous fix simply refused to tile whenever there was more than one stream. Give the reduce a separate dst_ne2 and use it for the destination index only. Ordinary FA split-K passes ne2 for it and is bit-identical. The sparse path passes the full batch, so tiling is unconditional again. This matters for memory, which is the binding constraint on this model. The scratch is capped at 256 queries per tile rather than scaling with the batch, so at ub2048 a 2-sequence context drops from 1028 MB to 128.5 MB and a 4-sequence one from 2056 MB to 257 MB. It also removes a cliff: 8 sequences at ub2048 would have exceeded maxStorageBufferRange and silently fallen back to the slower non-split kernel. Sparse FA timings are unchanged (every shape within 0.45%, inside run-to-run spread). 13,308 FLASH_ATTN_EXT cases pass, which covers the ordinary split-K path that shares this shader. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Assisted-by: Claude (Opus 5)
Assisted-by: Codex
The gap and its consequence were diagnosed by Jaap Buurman: the sparse prefill path gates on batch >= 64 and gather-to-compact gated on batch == 1, so batch 2..63 fell through to dense attention over the whole compressed KV, at a cost that grows with context. That is where a speculative draft lands (n_max 3-5), which is why token generation dropped off sharply with DSpark enabled. This implements the fix for that diagnosis. Give each query token its own gathered top-k block rather than deduplicating into a union. No dedup pass, no atomics, and the size is bounded by n_kv_raw + n_batch*n_top_k regardless of depth. Cross-token rows are neutralised through the mask, which already encodes each token's selection: token t reads -inf on any block that is not its own, so the softmax cannot double count. The compact mask is token-major [n_batch][kv_c], which is what the GQA mask path already expects (m_stride 0, rows stepped by gqa_iq1 * m_row_len). Measured on gfx1151, test-backend-ops perf, medians of 2 launches, n_kv_raw=2304 n_top_k=512. Gather cost is flat in depth (894 us at batch 4 at every depth); dense is not: kv rows batch dense gathered speedup 11008 2 2194 us 680 us 3.23x 11008 4 2204 us 895 us 2.46x 11008 8 2217 us 2217 us 1.00x (gate declines: kv < 2*kv_c) 35584 2 7065 us 680 us 10.38x 35584 4 7099 us 895 us 7.93x 35584 8 7114 us 1316 us 5.41x 133888 2 7552 us 680 us 11.10x 133888 4 7971 us 894 us 8.92x 133888 8 8348 us 1320 us 6.32x Batch 1 is unchanged in behaviour and stays on the same code path. Not validated end to end. A deduplicated union would shrink the gathered set further wherever adjacent draft tokens select overlapping keys, and would lift the batch ceiling documented in the following commit. Suggested-by: Jaap Buurman (@Mushoz) Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Assisted-by: Claude (Opus 5)
Both additions sit close to a limit that is invisible at the call site. The parallel Lightning Indexer uses 62720 of 65536 bytes of shared memory at N_WAVES=8 / HEADS_PER_TILE=4, so exactly one workgroup fits per CU. That is the intended trade, but raising either constant overruns the budget and the pipeline then fails to create and silently falls back to the small variant. Write the arithmetic down next to the arrays. The small-batch gather's compact set is independent of context depth but grows with batch, so its attention work is quadratic in batch against dense's linear. The kv >= 2*kv_c gate already caps this at about batch 6 at 32k depth and ~30 at 128k, and beyond the cap dense runs instead, so it is never slower. But the measured speedups do not show that ceiling and it is the main argument for building the deduplicated union later. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Assisted-by: Claude (Opus 5)
Replaces the per-token top-k blocks with one row per DISTINCT selected key, so the compact set stops growing linearly with batch. Measured on the real model, adjacent tokens share 60% of their selections over 4 tokens and 76% over 8, so the union is materially smaller than n_batch*n_top_k. The size is only known on the GPU, so flash-attention now reads its KV bound from a buffer instead of the push constant, behind a new DYNAMIC_KV pipeline flag. Every other pipeline folds the flag away at compile time and is unchanged (13308 FLASH_ATTN_EXT cases pass identically with the union off). No indirect dispatch is needed: FA workgroup counts come from neq1/neq2/neq3 and never from KV, so only the loop bound moves. Padding the count to 256 keeps KV % Bc == 0, which lets the aligned pipeline variant still apply. Dedup marks a bitmap from the top-k lists and compacts by scanning bitmap WORDS. An earlier version scanned the mask row by row: simpler, but it made dedup cost scale with depth and measured 0.71x at kv=133888, ie a loss. The bitmap form is O(n_batch*n_top_k) to mark and R/32 to compact, and is depth-independent. test-backend-ops perf, medians of 2, n_kv_raw=2304 n_top_k=512, fixture overlap 60% (the default generator produces near-zero overlap and would make the union look worthless by construction): kv batch per-token union speedup 35584 2 680 us 632 us 1.08x 35584 4 892 us 720 us 1.24x 35584 8 1315 us 900 us 1.46x 133888 2 679 us 638 us 1.06x 133888 4 892 us 726 us 1.23x 133888 8 1314 us 909 us 1.45x Union cost is flat in depth (720 vs 726 us at batch 4 across a 3.8x depth range), which the per-token form was not. Opt-in via GGML_VK_FA_TOPK_UNION=1, single stream only, and it falls back to the per-token blocks when the compressed region exceeds the shared bitmap. The kv >= 2*kv_c gate still uses the worst case, so the batch ceiling is NOT yet lifted; doing that needs a GPU-side fallback for an oversized union. Suggested-by: Jaap Buurman (@Mushoz) Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Assisted-by: Claude (Opus 5)
…size
The union shrinks the compact set but the gate could not see it. kv_c is the worst
case n_kv_raw + n_batch*n_top_k, so at batch 8 and 32k depth the gate priced 6400
rows against 11008 source rows, declined, and dense attention ran over the whole
compressed KV -- even though those selections deduplicate to about 3300 rows. That
is the batch ceiling the previous commit documented and did not lift.
The union size is only known on the device, and the note on the previous commit
assumed lifting the ceiling therefore needed a GPU-side fallback for an oversized
union. It does not. The union is bounded by the source, so there is no dispatch to
recover from; the missing piece was only a measurement. The union shader now writes
its count into a small host-visible buffer and the host prices the next step from
it. The read is deliberately unsynchronised and one graph stale -- overlap is a
property of the model and the draft, not of one op -- and a wrong estimate costs
part of one step, never correctness, because every allocation and dispatch bound is
still the worst case.
Estimates are held per batch size. A speculative decode varies the batch with the
accept count, and the overlap itself varies with the batch (0.64 at 2 tokens, 0.40
at 4, 0.24 at 8), so a single slot would be invalidated on nearly every step.
Each estimate tracks the latest measurement, lightly smoothed, and a declined step
dispatches the scan in a new count_only mode that writes no index list. Both of
those are deliberate and were measured the other way round first. Holding a decaying
peak to stay conservative, and sampling the probe every 256 declines to stay cheap,
together produced 1992 us on a shape the union runs in 898: the asymmetry actually
runs the other way, because an estimate that is too high declines compaction and
forgoes 2-3x for as long as it stays high, while one that is too low costs a single
step and is corrected by the count that step produces. Sampling makes it worse still,
since a decline is exactly the state in which the compact path stops refreshing the
estimate, so the stale value stays latched for the whole sampling period. The probe
is one workgroup against the ~2.2 ms dense op it rides along with, and it costs 0.2%
of a declined step.
The same measurement settles the reverse case. Where selections do not overlap, the
union is the same size as the per-token blocks and its scan is pure cost, so those
now stay on the per-token path rather than being admitted whenever the worst-case
gate happened to allow them.
test-backend-ops perf, medians of 3 launches counterbalanced A B B A A B, spreads at
or under 0.4%, n_kv_raw=2304 n_top_k=512, at kv=11008 (~32k source) which is the
shape where the gate used to decline. The fixture's ov is a per-token share, not the
union/selected ratio the model was measured by: at nb tokens it yields
(ov + (1-ov)*nb)/nb of the selections, so ov=86 is the setting that reproduces the
0.243 measured on DeepSeek-V4-Flash over 8 adjacent tokens, and ov=60 is deliberately
more pessimistic than the model.
batch overlap before after
8 ov=86 2215.2 us 704.8 us 3.14x
16 ov=86 2227.3 us 841.3 us 2.65x
8 ov=60 2214.5 us 897.9 us 2.47x
16 ov=60 2226.2 us 2231.2 us 1.00x union does not fit, declined
8 ov=0 2213.5 us 2219.0 us 1.00x nothing to deduplicate, declined
The union cost is depth-independent as before, so the lifted cells now sit alongside
the deeper ones: 897.9 / 900.7 / 906.6 us at batch 8 ov=60 across kv 11008 / 35584 /
133888. Every cell at kv 35584 and 133888 is 1.00x: this changes which shapes are
admitted, not how the union performs once admitted.
Both declines above are the estimate working rather than failing. At ov=60 and batch
16 the union really is 5888 rows against 11008 source rows and compaction would not
pay; at ov=0 there is no overlap to exploit at all. Verified from the executed graph
with GGML_VK_FA_UNION_STATS=1, added here, which reports the measured union/candidate
ratio and the resulting decision rather than leaving engagement to be inferred from a
timing. It reads 0.246 at ov=86 and batch 8, against the 0.243 measured on the model.
13310 FLASH_ATTN_EXT cases pass with the union off and on, including three new
overlap cases at the shapes the gate now admits.
End to end on a model that has no top-k tensor at all, so the compaction path is never
reached and the change should be structurally inert: Qwen3-Coder-30B-A3B UD-Q4_K_XL,
same counterbalanced order, medians of 3. pp512 1578.20 -> 1591.67 t/s and tg32 96.89
-> 97.28 at d0; pp512 700.64 -> 699.32 and tg32 53.92 -> 54.45 at d16384. All within
1%. The union path itself cannot be checked this way: it needs batch 2..63 decode,
which only a speculative draft produces.
Still measured against a synthetic fixture rather than real draft tokens: that needs
a runnable target and draft pair, which this box cannot host.
Suggested-by: Jaap Buurman (@Mushoz)
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Assisted-by: Claude (Opus 5)
…xture calibration The previous commit calibrated its fixture against a PROXY, because a proxy was all that was reachable at the time: adjacent PREFILL tokens standing in for draft tokens. With the draft now drivable, the proxy turns out to have been optimistic, and the correction runs against that commit's own message. DeepSeek-V4-Flash UD-IQ3_XXS with the DSpark draft, Vulkan attention and CPU experts, 25538 tokens of context. The gate reports union/candidate 0.560 and 0.609 on the target layers and 0.484 to 0.651 on a second shape (n_kv_raw 256, most consistent with the draft's own sparse attention, also GPU-resident). Mean about 0.56 at batch 4, against 0.397 from the proxy at a window of 4. Draft tokens diverge from each other more than adjacent prefill tokens do, which is what a draft exploring rather than committed text should look like. The fixture's ov maps to (ov + (1-ov)*nb)/nb, so ov=60 predicts 0.55 at batch 4 against 0.56 measured: ov=60 reproduces real drafting almost exactly. The previous commit says ov=86 is the setting that reproduces the model and ov=60 is deliberately pessimistic. That is backwards for real draft tokens, and the realistic headline at batch 8 is 2.47x rather than 3.14x. The kernel numbers in that table are unchanged and still correct for the overlap each column states; only which column describes reality has moved. Batch 8 overlap itself is still not measured on real drafts, for the reason below, so 2.47x is an extrapolation from batch 4. The path does engage: 56 of 57 sampled gate decisions took the union, and the single decline is the unseeded first call, which then probes. That is the designed bootstrap, observed on the real model rather than a fixture. Scope, stated plainly because the previous commit implies more: DSpark drafts 3 tokens, so the batch is 4 and --spec-draft-n-max does not raise it, an MTP head having a fixed width. At batch 4 and this depth kv_c is 4352 against n_kv 8704, so the OLD worst-case gate admits too and the gate change makes no difference there; the smaller compact set comes from the union commit. For this target and draft the gate change buys a lower depth threshold instead: roughly 21.0k tokens rather than 25.5k on the target layers, and 11.8k rather than 17.7k on the second shape. The 2.47x and 2.65x cells need batch 8 or 16, which this pair does not generate. GGML_VK_FA_UNION_STATS takes a period now rather than being fixed at 256, and reports running totals. The fixed period is why this took two runs: a 64-token speculative decode over 25k of context emitted exactly ONE line, which reported only that the first call was unseeded. A diagnostic whose sampling rate can hide the thing it exists to measure is not a diagnostic. Measurements and provenance: ~/strix-results/derived-dsv4-union-gate-20260814.md section 9. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Assisted-by: Claude (Opus 5)
Every other gate in this family is on by default with =0 to disable: GGML_VK_FA_TOPK, GGML_VK_FA_TOPK_CM, GGML_VK_FA_TOPK_SPLIT. GGML_VK_FA_TOPK_UNION=1 was the odd one out because it began as a prototype, and it is no longer one. What makes it safe to default is that the path declines itself rather than needing a user to know when to avoid it. Where the selections do not deduplicate, the estimate says so and the per-token form runs instead; where the compact set would not fit under the gate, dense runs. Both were measured, at 1.00x, rather than assumed. Kept as one commit of its own rather than folded into the gate change, because it is a policy change and not a mechanism one: reverting these four lines restores opt-in behaviour without touching the gate logic, which is what a bisect would want if a regression turns up. 13310 FLASH_ATTN_EXT cases pass with the variable UNSET, which is the configuration this changes and the one that ships. That is the same code path as the previously validated GGML_VK_FA_TOPK_UNION=1, but it was run again rather than argued from equivalence. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Assisted-by: Claude (Opus 5)
The gather rejected anything but f16 K/V, so a DSv4 run with -ctk q8_0 -ctv q8_0 took the
dense fallback for every batch 2..63 decode - the whole small-batch path, silently off. Two
community test sets were run that way before anyone noticed, which is what prompted this.
The restriction was incidental rather than essential. The gather RELOCATES rows; it never
reads a value out of one. Addressing K as raw 4-byte words instead of f16 elements makes the
same shader serve any type whose row is a whole number of words, which every block-quantised
type here satisfies at head size 512. The union shader needed no change at all - it only ever
touched top_k indices.
Zeroing an unused row now writes zero BYTES. That decodes to zero for the block-quantised
types as well as for f16, because a zero scale zeroes the block; those rows are -inf in the
mask regardless, so zeroing only keeps a garbage dot product out of the softmax as a NaN.
The type gate asks ggml_vk_fa_kv_native rather than just "is it quantised". While the compact
scratch is active the dequant/contiguize pass is disabled by construction, and the assert
guarding that combination aborts rather than falling back, so admitting a type flash-attention
has no native shader for would crash instead of degrade.
test-backend-ops perf, kv=11008 n_kv_raw=2304 n_top_k=512 ov=60, gather off vs on:
batch f16 dense -> gathered q8_0 dense -> gathered
2 2192 -> 645 us 3.40x 3917 -> 1110 us 3.53x
4 2203 -> 726 us 3.03x 3924 -> 1257 us 3.12x
8 2213 -> 910 us 2.43x 3936 -> 1567 us 2.51x
16 2219 -> 2232 us declined 3930 -> 3952 us declined
Two things worth stating because they contradict what I expected going in.
Compaction does NOT pay more on quantised K/V. The reasoning was that a gathered set is
contiguous and would dodge the strided-read and channel-aliasing taxes that hurt quantised
reads more; the measured ratios differ by 3 to 4%, inside noise. It pays the same proportion
from a worse starting point.
That worse starting point is the real result: q8_0 attention is about 1.78x SLOWER than f16
here, dense (3917 vs 2192) and gathered (1110 vs 645) alike. Halving the bytes does not pay
for the dequant work in the inner loop at this head size. So this change makes q8_0 K/V
survivable for DSv4 rather than advisable - f16 is still the faster cache by a wide margin,
and that inverts the usual q8_0-KV-buys-speed rule of thumb for this model.
13318 FLASH_ATTN_EXT cases pass, including 8 new q8_0/q4_0 top-k cases at the shapes a DSv4
decode actually hits. test_flash_attn_ext_top_k takes a type_K parameter for them.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Assisted-by: Claude (Opus 5)
Relocating quantised rows verbatim leaves flash attention to decode them in its inner loop,
which it does once per query block that reads a row rather than once per row. Measured, that
costs 0.15 us per KV row attended - 0.1481 / 0.1510 / 0.1512 us/row at batch 8 / 4 / 2 and
0.1565 dense, so a constant across a 3.6x span of row counts and across both regimes. It is
what made q8_0 attention 1.78x slower than f16 while reading half the bytes.
The gather is the natural place to pay it instead: it already touches every selected row
exactly once. Decoding there makes the compact scratch f16, flash attention takes its f16
path, and the redundancy disappears. No extra pass - the pass already existed.
test-backend-ops perf, kv=11008 n_kv_raw=2304 n_top_k=512 ov=60, q8_0 K/V:
batch dense verbatim gather decoded gather f16 gather
2 3917 us 1110 us 630 us 640 us
4 3924 us 1257 us 715 us 727 us
8 3936 us 1567 us 891 us 908 us
q8_0 now edges f16 by 1.5 to 1.8% rather than trailing it by 78%: the FA does identical f16
work either way, and the gather's scattered read side moves half the bytes. Against the dense
fallback a q8_0 cache took before any of this work, batch 8 is 4.41x.
Flash attention picks its pipeline from what the SCRATCH holds, not from the source tensor, so
k_type_eff/v_type_eff now also key off the compact state - the same mechanism use_dequant_kv
already used, which is why that variable existed in this form.
q8_0 only, deliberately. It is the type both community test sets used and the one worth having;
every other quantised type keeps the verbatim path, which is correct for them and validated.
Generalising means one shader variant per type through dequant_funcs.glsl, which is only worth
the shader-permutation cost if this shape of win reproduces elsewhere.
Not covered: the per-token gather still relocates verbatim, so a batch where the union declines
but the per-token form fits keeps paying the inline decode. Same fix applies, and the same
measurement will say whether it is worth a second shader.
13318 FLASH_ATTN_EXT cases pass.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Assisted-by: Claude (Opus 5)
q4_0 turns out to want this more than q8_0 did. Its inline-decode penalty is the same shape - a constant per KV row attended, flat across batch and across dense vs gathered - but larger: 0.172 to 0.179 us/row against q8_0's 0.148 to 0.153. Unpacking nibbles costs more ALU than the halved bytes save, which is also why q4_0 was the SLOWEST of the three types before this. kv=11008 n_kv_raw=2304 n_top_k=512 ov=60, batch 2 / 4 / 8 type dense verbatim gather decoded gather f16 2192/2203/2213 - 640/726/902 q8_0 3917/3924/3936 1110/1257/1567 628/712/888 q4_0 4134/4141/4145 1189/1342/1671 626/711/886 All three converge, because after the gather they run the same f16 attention; what is left is the scattered-read side, where fewer bytes now wins. Against the dense fallback each type took before any of this work, batch 8 is 2.45x for f16, 4.43x for q8_0 and 4.68x for q4_0. The interesting part is what the first attempt got wrong. Generalising via dequantize4() from dequant_funcs.glsl looked obvious and passed on q8_0, then failed q4_0 with NMSE 1.17. That helper is permutation AGNOSTIC by design: mul_mat_vec only ever feeds it into a dot product, where any consistent element order gives the same sum, so for q4_0 it returns nibbles in packed order rather than element order. q8_0 passed only because its layout happens to be contiguous. Materialising a row into memory is precisely the use where the order matters, so each type's true mapping is now written out - for q4_0, byte j carries element j in its low nibble and element j+16 in its high one, exactly as dequant_q4_0.comp materialises it. That is also why this stops at two types rather than looping over the whole native list: every type needs its mapping stated and tested, an #error guards anything else reaching the shader, and q8_0 and q4_0 are the two actually used as a KV cache. q4_1/q5_0/q5_1/iq4_nl keep the verbatim gather, which is correct for them. 13318 FLASH_ATTN_EXT cases pass, including q4_0 and q8_0 top-k cases at the shapes a DSv4 decode hits, plus q4_0 perf cells alongside the q8_0 ones. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Assisted-by: Claude (Opus 5)
Generalising the gather from f16 elements to raw words (6b2cade) replaced [[unroll]] for (uint i = 0; i < HEAD_SIZE / LANES; ++i) with a loop bounded by the row_words push constant, which cannot be unrolled and pays a bounds check per iteration. Half as many, twice as wide accesses were still 6% slower on the smallest op the gather serves. Bisected over the 19 beta3 commits with the tip's test file overlaid at every point, so the instrument stayed fixed while beta3 grew its own perf cases. nb=1 at kv=8192: 54.58 us at the pre-beta3 base, flat through commit 16 (54.84), 58.24 at 6b2cade, 58.23 at the tip. Step 4*LANES instead. Each access stays contiguous across the lanes, so coalescing is unchanged, and an f16 row at head size 512 is 256 words == 4*LANES: one iteration, no loop overhead. The tail loop carries the quantised row sizes, which are not multiples of 4*LANES (q8_0 is 136 words, q4_0 is 72). nb=1 against the pre-beta3 base, f16 K/V: -0.07% at kv=8192, -2.83% at kv=35584, mean -1.29% across the six shapes, so the regression is gone rather than reduced. Against beta3 as merged: -4.8% to -6.2%. The nb>=2 gains are untouched, worst cell -0.51%. 13318 FLASH_ATTN_EXT cases pass. Worth noting the gather itself was never the cost at nb=1: with GGML_VK_FA_TOPK_GATHER=0 that cell is 249.84 us against 58.14 us with it on. Assisted-by: Claude Opus 5
The union gather decodes q8_0/q4_0 on the way in, but it exists only for n_batch > 1: a single token has nothing to deduplicate, so plain autoregressive decode takes the per-token gather, which relocated rows verbatim and left flash attention paying the same inline decode the union work measured at ~0.15 us per KV row attended. flash_attn_gather_dq.comp closes that: same row mapping as flash_attn_gather.comp, same per-type element mapping as flash_attn_gather_union_dq.comp (q4_0 materialised in element order, not packed order). The compact scratch becomes f16 and flash attention takes its f16 path, keyed off the compact state exactly as the union path already is. test-backend-ops perf, kv=11008 n_kv_raw=2304 n_top_k=512 ov=60, nb=1: type verbatim gather decoded gather q8_0 226.9 us 102.9 us q4_0 258.4 us 101.8 us f16 114.5 us 105.5 us (f16 never decodes; its two runs bound the harness spread) Both quantised types land on the f16 op time instead of trailing it 2.0-2.3x: after the gather the attention work is identical f16, and the gather's scattered read side moves fewer bytes. q8_0 and q4_0 only, as in the union: each type's element mapping must be stated and tested, an #error guards anything else reaching the shader, and every other type keeps the verbatim gather. nb=1 rows are added to the top-k perf grids so this regime stays measured. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Assisted-by: Claude (Opus 5)
ggml_vk_flash_attn_top_k declined any K/V but f16, so a DSv4 run with a quantised cache took dense FA for every prefill batch - O(kv) against the sparse path's O(n_kv_raw + n_top_k), a penalty that grows with depth. A 128GB community run measured it end to end: prefill -5.4/-22.2/-39.4/-52.9% against f16 at source depths 17k/33k/67k/134k with -ctk q8_0. The sparse shaders stay f16-only. A quantised cache is dequantised once per op into the prealloc_x f16 scratch by the dense path's own fused dequant+transpose pipeline, and the shaders read the scratch. K and V are the same tensor in this path, so one pass covers both at half the dense path's scratch footprint. Types without a dequant-transpose pipeline, oversized caches, and GGML_VK_FA_DEQUANT=0 keep the dense fallback. test-backend-ops perf at nb=1024, kv 5504/11008/19200/35584 (the community depths in compressed-K rows): q8_0 was 60.9 -> 427.6 ms dense (1.30x -> 8.27x of f16, growing with kv) and lands within +0.8% of f16's flat 46.9-51.7 ms with the scratch; the dequant pass itself costs under 1%. 36/36 FLASH_ATTN_EXT top-k cases pass, including 6 new quantised prefill cases (nb=64 ns=1/2, nb=128 split form). Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Assisted-by: Claude (Opus 5)
The K tile is invariant across the head loop but coopMatLoad sits inside it, so the shader re-read the same 8 MatrixA fragments from shared memory once per head tile - 4x at N_HEAD=64. Each wave64 fragment load moves 1 KiB of LDS traffic because the 16x16 f16 fragment is replicated 4x across the subgroup. Same loads of the same data feeding the same coopMatMulAdd sequence, so the output is bit-identical; the fragments just live in registers (8 of them, 16 f16 per lane) instead of being re-fetched. gfx1151, test-backend-ops perf, kv 8704/33280/131584 x batch 1-15: 16-20% faster at every shape, no spills, no regression. Batch 1 is ordinary DSv4 decode and gets 12-18% of that. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Assisted-by: Claude (Opus 5)
…ecode CM shader Builds directly on pepuscz's PR halo-box#6 and issue halo-box#10. Their per-kernel table - 43.8 us scalar against 27.9 us small CM per 1k scanned rows at batch 5, and the same ratio at every depth - is what showed the indexer's cost is a per-tile constant rather than anything to do with tokens or bytes. Once the cost is per tile, the thing to minimise is the tile count, and that is what this changes. The finding is downstream of their work; only the choice of shader differs. The decode CM shader puts 16 HEADS in the coopmat N dimension and dispatches one workgroup per token, so it issues 4*n_batch tiles per 16 KV rows. That is the arithmetic minimum - (64 heads x n_batch tokens) / 16 columns - because 64 heads fill its 16 columns exactly, with no remainder at any batch. The small CM shader puts 16 TOKENS in N and pays a flat 64 tiles however small the batch is, so at batch 5 eleven of its sixteen columns are padding. Measured cost tracks tile count: at kv=131584 batch 1 costs 1.95 us per 1k scanned rows and batch 5 costs 26.28, i.e. 13.5x for 5x the work, and both shaders sit at 6.6-7.8 ns per tile. The decode CM shader body has no n_batch == 1 assumption - token is gl_WorkGroupID.y and it indexes q/w/mask/dst by it - so the old gate was an artefact of where it was written. gfx1151, kv=131584 (526k source tokens), us/run, shipped vs this: batch 2 2335.6 -> 419.0 (5.6x, was on the scalar path) batch 3 3492.3 -> 631.8 (5.5x, was on the scalar path) batch 4 3267.9 -> 840.8 (3.9x) batch 5 3465.6 -> 1119.4 (3.1x) <- DSpark n-max 4 verify shape batch 8 3704.9 -> 1764.1 (2.1x) batch 15 4467.9 -> 3231.6 (1.4x) Batch 16 and 32 are unchanged, which confirms the arms are isolated. PR halo-box#6's route is kept as the opt-out arm rather than deleted, and is promoted from opt-in to default-on so that one variable is enough to reach it: default decode CM for the whole 2-15 window ..._DECODE_CM_BATCH=0 small CM for 4-15, scalar for 2-3 (PR halo-box#6) ..._DECODE_CM_BATCH=0 SMALL_CM=0 scalar for 2-15 (pre-PR halo-box#6 baseline) Verified all three route as documented and pass 29/29. Keeping it is not just courtesy: the tile-count argument is hardware independent, but decode CM re-reads the K tile once per token and that part is bandwidth dependent, so the crossover need not sit in the same place on other devices, and these numbers are from one gfx1151 box. A (head, token) packing shader was also tried and is strictly worse - it reaches 4b tiles only when the batch divides 16, and it divides the workgroup count by the batch. Kept out of tree. Numerics: decode CM sums the 64 heads in tiles of 16 rather than one at a time, so f32 accumulation order differs from the small CM path. Cleared by KLD A/B on trunc10 (-c 8192, -b/-ub small so every indexer dispatch routes through the window, 8 chunks wikitext-2, same binary, only the env var differing): -ub 8 mean KLD 0.000000, max 6.0e-5, RMS dp 0.000%, same-top 99.994%, PPL 10507558.3156 identical in both arms -ub 5 mean KLD 0.000000, max 8.8e-4 but 99.9% 4.9e-5 (one tail event, no argmax flip), RMS dp 0.000%, same-top 100.000% Same class as PR halo-box#6's own measured numbers (max 5.5e-5, same-top 100.000%) and 45x tighter than the FA_WAVE32 change already shipped. Adds the eval coverage the 2-15 window never had (batches 4/8/15 at kv=256) and the perf grid these numbers came from. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com> Assisted-by: Claude (Opus 5)
Same defect as the decode CM shader, larger blast radius. `k_sh[wave]` is written
once above the head loop and never touched again, but `coopMatLoad` sits inside
both the head_base and head_local loops, so each wave re-reads the same 8 MatrixA
fragments N_HEAD times: 512 loads where 8 would do. Each wave64 fragment load
moves 1 KiB of shared-memory traffic because the 16x16 f16 fragment is replicated
4x across the subgroup.
Same loads of the same data feeding the same coopMatMulAdd order, so the output is
bit-identical; the fragments live in registers (8 of them, 16 f16 per lane) instead
of being re-fetched.
This source compiles to both the prefill pipeline and the small-batch one, so both
gain. gfx1151, test-backend-ops perf, us/run:
prefill (lightning_indexer_cm_f16)
kv 8704 nb 2048 27492.5 -> 18207.0 1.51x
kv 33280 nb 2048 104101.9 -> 68947.8 1.51x
kv 131584 nb 2048 421239.8 -> 281540.5 1.50x
kv 131584 nb 16 3456.2 -> 2329.2 1.48x
kv 131584 nb 32 6858.6 -> 4623.6 1.48x
small batch (lightning_indexer_cm_small_f16)
kv 131584 nb 5 3433.1 -> 2783.3 1.23x
kv 131584 nb 8 3682.3 -> 3105.1 1.19x
kv 131584 nb 15 4474.4 -> 4209.5 1.06x
1.48-1.54x on prefill at every depth and batch measured, flat. The small-batch
variant gains less, which says the 1-wave configuration is bound by its Q staging
and epilogue rather than by fragment traffic.
No regression at any shape, so the 8 held fragments (64 VGPRs per lane) neither
spill nor cost occupancy even in the prefill pipeline, which runs 512 threads at
62.7 KiB of shared memory and was already pinned to one workgroup per CU.
LIGHTNING_INDEXER 29/29.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Assisted-by: Claude (Opus 5)
…ation constant
The indexer pipelines hardcoded nh=64 (DeepSeek-V4's geometry) in supports_op and
as a shader constant, so any other head count fell to the CPU backend.
Qwen3.8-Flash-Next (qwen4_exp, released today) carries the same indexer at nh=4.
N_HEAD becomes specialization constant 1, one pipeline per supported head count
(LI_NH_VALUES = {4, 32, 64}), selected by q->ne[1] at dispatch. It must be a
specialization constant, not a push constant: the head loop's trip count has to
stay visible to the compiler - as a push constant the loop cannot unroll and
decode costs ~55% (measured, gfx1151, kv=131584, and why this commit is not the
simpler design).
nh=4 does not fill a 16-wide head tile, so the tail tile stages q as zeros
(relu(0)*w = 0, padded heads drop out of the sum) and the weight read clamps its
index. Both guards are gated on N_HEAD % TILE != 0, a specialization-constant
expression, so at nh=64 they fold away at pipeline compile and the codegen is
equivalent to the previous shader - measured, because relying on the compiler to
range-prove head < N_HEAD instead still cost ~45%:
kv=131584, us/run before after
decode b1 226.8 219.6
decode b5 1058.8 1007.4
decode b15 3174.0 2966.8
prefill b16 2329.2 2230.2
prefill b2048 281540.5 281169.1
No regression at nh=64; the small decode improvement is within a cautious
reading of run-to-run variance and is not claimed.
test-backend-ops LIGHTNING_INDEXER: 44/44, up from 29/29 - the suite's existing
nh=32 cases had been silently reporting "not supported" under the old gate and
now run against the CPU reference.
Head size stays pinned at 128: the coopmat tiling and f16vec4 staging both
assume it, and every known indexer model (DSv4, GLM-DSA, Qwen4-exp) uses 128.
Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Assisted-by: Claude (Opus 5)
The scalar64 shader kept its hardcoded 64-head loop when the pipelines went per-head-count; nh 4 and 32 would have run with the wrong trip count. Assisted-by: Claude Fable 5
The GEMM2 P load depends only on bc_chunk, but sat inside the hsv_tile loop, so all four fragments were re-read from shared memory once per tile (2 tiles at HSV=128, 4 at HSV=256). Load them once into a coopmat array before the loop. Psh is not written again until the next KV block, so the fragments stay valid across the barriers inside it. Measured on gfx1151 (RADV), Qwen3-Coder-30B-A3B, pp2048 f16 KV: d8192 +6.9%, d16384 +8.1%, d32768 +9.2%. VGPRs and LDS unchanged, 0 spilled. FLASH_ATTN_EXT suite 5105/5105. Assisted-by: Claude Fable 5
Psh held P as [kv][query]. The GEMM2 UseA load therefore had to request ColumnMajor, and RADV only attaches an alignment hint to the internal column-major case, which for UseA means the RowMajor request. The load was emitted as 16 separate 16-bit shared reads per fragment. Store P as [query][kv] instead and request RowMajor. The producer now writes four scalar components rather than one vec4; those go to disjoint bytes, so there is no read-modify-write and no partially written vec4. Also updates the host shared-memory estimator, which mirrored the old stride and would otherwise disagree with the shader. Perf-neutral on gfx1151 (within 1% at hd128 and hd256), but it shrinks Psh: LDS 16384 -> 15360 B and code size 13896 -> 13596 at hd128, VGPRs unchanged at 96 with 0 spilled. Kept because it is free and removes a scalar shared-memory access pattern. FLASH_ATTN_EXT suite 5105/5105. Assisted-by: Claude Fable 5
get_fa_tuning_params_coopmat1 took device->subgroup_size unconditionally, so on a 64-wide device the coopmat1 FA path always ran wave64. The sibling scalar path already does AMD-specific wave selection; the coopmat1 path never got the equivalent. Narrowing is free exactly when it does not add an iteration to the O-accumulation loop, which runs ceil((HSV/4) / threads_per_rowgroup) per row. On a 64-wide device the test reduces to hsv <= 128. Above it the narrow subgroup issues 1.5x to 1.8x the instructions for the same SIMD passes and hd256 measures 6 to 18 percent slower, so the rule declines. Gated behind GGML_VK_FA_WAVE32 (=1 rule, =2 forces the pin regardless of head size, diagnostic only). Off by default. Measured on gfx1151 (RADV), model-level pp2048, Qwen3-Coder-30B-A3B: d0 +2.5%, d8192 +8.4%, d16384 +10.1%, d32768 +11.3%. Op-level across head sizes, largest where wave64 wastes the most lanes: hsv=64 +12.3%, hsv=96 +7.4%, hsv=128 +6.5%, hsv=256 correctly declined. Full test-backend-ops suite 15884/15884 at =0, =1 and =2; the pin was confirmed to engage in-band via pipeline VGPR statistics rather than inferred from the env var being set. Assisted-by: Claude Fable 5
Split from the fork's flag-flip commit (f7d804ee7): GGML_VK_FA_WAVE32 now defaults to 1 (apply the HSV-based narrowing rule); =0 opts out and =2 keeps its diagnostic force meaning. The subgroup_size_control device guard is unchanged, so devices without a 32-wide subgroup mode are unaffected. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Assisted-by: Claude (Opus 5)
Select the matmul_id tile by expected per-expert token count (nei1*nei0/n_expert) instead of aggregate nei1, which always picked the widest tile and left most N-lanes empty at MoE prefill (~16 rows/expert vs BN=64). Gated on GGML_VK_MMID_SMALLN=1, default off pending cross-model validation. Was null (+0.7%) standalone pre-row-lists: smaller tiles meant more workgroups each re-paying the per-WG id scan. With the Stage 1 row-list prepass the scan is gone and the occupancy win materializes. Strix Halo (RADV gfx1151), Qwen3.6-35B-A3B UD-Q5_K_XL, clean window, pp512 4-way (rowlists x smalln): 914.7 / 937.9 / 1005.3 / 1063.6 t/s (combined +16.3% vs baseline). MUL_MAT_ID q5_K 2.33->4.43 TFLOPS (1.91x), q6_K 1.99->3.57 (1.79x) cumulative. Hot pipelines verified via probe: matmul_id_*_f32 (y_f32 path, quantize_y does not engage for these quants); smalln flips tile _m -> _s. 790/790 MUL_MAT_ID both rowlists configs. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Assisted-by: Claude (Opus 5)
Two env-gated overrides of the mmid small-tile config (KHR coopmat branch only, scoped to the mul_mat_id quant pipelines; dense untouched): - GGML_VK_MMID_TILE16: BN/WN 32->16. NEGATIVE (-3.8% e2e): experts with n_e>16 split into two column tiles and re-stream their full weight matrix; A-traffic scales with sum(ceil(n_e/BN)), so BN must not drop below the mean per-expert n. Kept as a documented dead end. - GGML_VK_MMID_BM64: BM 32->64, BLOCK_SIZE 64->128 (two warps). +1.3% e2e on top of rowlists+smalln: same A-traffic, half the ir-tiles so half the per-expert B re-reads, larger WGs hide latency. Strix Halo, Qwen3.6-35B-A3B UD-Q5_K_XL pp512, clean window, all on top of RL+SN control 1065.4 +/- 3.3: BM64 1078.9 +/- 3.5, TILE16 1024.6, BM64+TILE16 1007.5. Cumulative vs pre-Stage-1 baseline: 914.7 -> 1078.9 (+17.9%). 790/790 MUL_MAT_ID for all configs. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Assisted-by: Claude (Opus 5)
…ated) Extends the BM64 idea to the medium tile: GGML_VK_MMID_M128=1 raises the mmid m-tile to BM=128 / BLOCK_SIZE=4*subgroup (four warps), same A-traffic, half the ir-tiles and per-expert B re-reads. The medium tile is what the per-expert-n heuristic selects at n~64 (e.g. ub2048 on 256-expert/top-8 models, or ub1024 at 128 experts). Strix Halo, Qwen3.6-35B-A3B UD-Q5_K_XL pp2048, drain-verified window: ub2048 stack 1039.9 +/- 1.1 -> +M128 1091.6 +/- 2.6 (+5.0%); inert in the s-tile regime (ub1024 1137.2 vs 1147.9 ref). ub1024 remains the throughput sweet spot for this model. 790/790 MUL_MAT_ID. Also measured this session, NOT kept: caching counts+row lists across the three expert matmuls of a layer (they share one ids tensor) was correctness-clean but perf-null (1066.3 vs 1070.9 baseline) - the prepass dispatches and barriers are already free on this queue; reverted rather than carry the invalidation surface. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Assisted-by: Claude (Opus 5)
Force required subgroup size 32 on the KHR-coopmat mmid quant pipelines (pipeline_dequant_mul_mat_mat_id[*] only; dense untouched). Hypothesis: RDNA3.5 WMMA is wave32-native, so RADV may lower KHR_coopmat better at wave32 than at the reported default of 64. Mechanics: the cm1 path in mul_mm.comp derives its warp grid from the real subgroup (warp_i = gl_SubgroupID, tiw = gl_SubgroupInvocationID) while NUM_WARPS = BLOCK_SIZE/WARP (spec constants) sizes coopmat_stage[] and ballots_sh[] and warp_r/warp_c assume NUM_WARPS == (BM/WM)*(BN/WN). Forcing sg32 with WARP=64 would over-run those shared arrays and leave warp_c outside the tile, so the gate (a) sets WARP=32 in shadowed copies of the s/m/l mmid warptiles (composes after the BM64/M128 shadows) and (b) halves WM (or WN, keeping WM>=TM, WN>=TN) until the doubled subgroup count exactly tiles BM x BN again, asserting both invariants. BLOCK_SIZE is kept, so workgroup shape, load loops and shmem match the wave64 stack; per-lane accumulator footprint is also unchanged (half the lanes per subgroup, half the (WM/TM)*(WN/TN) fragments). The required size is passed via a scoped CREATE_MM redefinition adding a trailing required_subgroup_size arg (fp16-branch pattern), gated on subgroup_size_control covering 32 since ggml_vk_create_pipeline_func silently drops the required size otherwise. Correctness (test-backend-ops -o MUL_MAT_ID -b Vulkan0): 790/790 plain, 790/790 WAVE32=1, 790/790 WAVE32+SMALLN+BM64. Perf, Qwen3.6-35B-A3B-UD-Q5_K_XL, fa=1 b/ub=512 ctk/ctv=q8_0 pp512 r=3, atomic window, canary clean: stack (SMALLN+BM64), wave64 default: 1076.61 +/- 3.71 t/s stack + WAVE32: 1106.73 +/- 2.52 t/s (+2.8%) Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Assisted-by: Claude (Opus 5)
The header claimed the workflow had never been executed. Run 37682939495 on 2026-10-07 did execute it: checkout, HIP build, device detection, test-backend-ops and artifact upload all worked and a DevLab runner was assigned with the expected labels. test-backend-ops reported 31003/31004 with one marginal MUL_MAT_ID q4_1 failure, which skipped perf and llama-bench on that run, so no throughput number exists yet. Assisted-by: Claude Opus 5
llama : load per-layer embedding tables from a separate GGUF (--ple)
…t-cache" This reverts commit 0f36c46, reversing changes made to 307c50d. --cache-dir drives runaway disk writes on ROCm: the slot state is copied straight between the GPU and a MAP_SHARED file, the runtime pins the mapped pages, and writeback keeps rewriting the file until the server exits. Users reported ~1.6 TB written to NVMe in an hour, and decode stalling under 5+ GB/s of continuous writes with a 75k-token context. Reproduced on gfx1151: one 503 MiB entry became 5.7 GB of writes within 60 s of idle. Even with that fixed (halo-box#149), an agent-style workload rewrites the whole state on every turn. Remove the feature until it can come back with bounded writes; --cache-ram is unaffected. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01G5NBdiBqRmGomhSRXBoucx
Revert "server : add --cache-disk for a persistent prompt cache" (halo-box#81)
SOLVE_TRI aborts under HIP graph capture: hipBLAS returns CUBLAS_STATUS_INTERNAL_ERROR from hipblasStrsmBatched inside solve_tri_f32_cublas (ggml-cuda/solve_tri.cu:71) and ggml aborts, so the perf step died after 10m23s having produced 138 rows. Reproduced locally at 4977288 on ROCm 10.0 and seen on the DevLab host on ROCm 7.2.1; test mode passes 24/24 on the same binary, so it is specific to the perf path under capture. The branch does not touch solve_tri.cu or the capture code. GGML_CUDA_DISABLE_GRAPHS is a runtime switch and lets all five SOLVE_TRI perf cases complete. It is set only on the perf step: llama-bench keeps graphs on because that is the production configuration. Remove it once the capture interaction is fixed. Assisted-by: Claude Opus 5
…spatch cuda : qwen4exp decode throughput on gfx1151 (small-row matvec dispatch, GDN rollback fusion, PR halo-box#106 ports)
spec : `--spec-n-rs-seq` caps the recurrent rollback snapshot slots; env var for `--spec-draft-mtp-vocab`
vulkan : WMMA selected-key attention for QSA block selection
…cratch flash_attn_cm1_mr.comp: one workgroup of eight wave32 subgroups owns 64 rows (H heads x 64/H tokens that share a KV head). Q stays in registers for the whole KV loop and every 16-key K/V tile serves all 64 rows, where cm1 pins Br at 16 and streams the whole K/V range per 16 rows. K is read as per-head f16 rows, V from a per-head V^T scratch written by dequant_f16_transpose_vt.comp ([HSV][KV + 128], 64x64 LDS tiles). Keys masked in every output row have their V zeroed per tile, so a WMMA -0.0 product cannot move the output and NaN V in masked cells cannot leak. GGML_VK_FA_MR: default on for head size 256 on RADV RDNA3, =1 for every eligible shape (head size 128 or 256, f16 K after the contiguize pass, f32 accumulation, KV % 64 == 0, mask shared across heads), =0 off. GGML_VK_FA_MR_LAZY / _NOEB / _VPOST (default on) are schedule knobs. GGML_VK_FA_VT: V^T for the cm1 kernel too (default on RADV RDNA3 when NAN_SAFE_V and shmem staging are off), =0 also disables MR. New flag bits 64..1024; NAN_SAFE_V keeps 32. Measured in the fork (v0.7.8) on gfx1151, Qwen3.8-27B UD-Q4_K_XL: FA per pp2048 ubatch at d32768 2067 -> 954 ms, pp2048 +21.5% at d32768, +5.7% at d8192, KLD 0.000197 against the cm1 kernel. Assisted-by: Claude Opus 5.5
…CONV + SILU) Gated DeltaNet layers build CONCAT(state, transpose(x)) -> SSM_CONV -> SILU, plus copies of the last 3 concat columns into the recurrent conv state (one per rollback slot). The concat transposes x into a [T+3, C] tensor that the conv reads once: 84 MB written and read per layer for Flash-Next at a 2048-token ubatch. This follows the CUDA backend's gdn_conv match (same graph, no llama change): - at the CONCAT, when its only readers are one SSM_CONV (+ SILU right after) and conv-state tail copies, write only the tail columns those copies read; - at that SSM_CONV, run a direct conv + SiLU: one thread per channel and 8 tokens, the 4-tap window sliding in registers, lanes along channels, the same dot() as the existing nc == 4 path (identical output). The decision is made once at the concat and consumed at its conv. The match checks the allocated addresses (the SiLU output overlaps nothing the conv reads, and nothing between the concat and the conv writes x or the state), and graph_optimize asks the allocator to keep x and the state alive through the SiLU. x and the state are also added to the barrier tracking. Ubatches under 8 tokens per sequence keep the concat (decode/verify unchanged). GGML_VK_SSMCONV_DIRECT=0 restores the concat + SSM_CONV_SILU pair; GGML_VK_FUSION_DEBUG=1 prints why a concat did not match. test-backend-ops: new CONCAT_SSM_CONV_SILU cases (scheduler-allocated, verify the SiLU output and every state copy): n_t 1..300, ragged channels and token tiles, two sequences, 1/4/8 rollback slots (8 reaches into the state columns), 10240 channels, and a whole-concat extra reader that must not fuse. Doubling the direct-conv store fails the fused cases + GDN_CONV_PREFILL; doubling the tail store fails the fused cases' state copies; both pass with GGML_VK_SSMCONV_DIRECT=0. Only tested on gfx1151 (RADV). llama-bench pp2048 -ub 2048, same binary with the gate off/on, ABBA x2, -r 2 (mean of 8): Flash-Next REAP-320 778.4 -> 790.3 t/s (+1.5%), Qwen3.6-35B-A3B 2101.8 -> 2161.0 (+2.8%). Every gate-on sample beat every gate-off sample. Flash-Next PPL c4096 (2 ubatches per chunk) identical per chunk. Assisted-by: Claude Opus 5.5
… fusions rms_norm.comp runs one 512-thread workgroup per row, so on 128-wide rows (the GDN head norms) 3/4 of the lanes idle and every row pays a shared-memory tree reduction. rms_norm_small.comp gives each row one subgroup (subgroupAdd), for ne00 <= 256, f32 in/out, contiguous rows. It handles plain RMS_NORM and RMS_NORM_MUL, plus two fusions only it has: - RMS_NORM_MUL_MUL: the gated norm, normalised * gamma * gate (qwen3.5/qwen4exp build_norm_gated). The gate chain (sigmoid/silu of z) sits between the two MULs in the graph, so graph_optimize now schedules it ahead of the norm and keeps the three nodes adjacent. The products are the same f32 multiplies as the separate ops. - RMS_NORM_SCALE: build_gdn_l2_norm, scale(rms_norm(x, eps/n), 1/sqrt(n)). The factor is applied on the store after rounding the normalised value to f32, the same two roundings as the SCALE op. graph_optimize keeps the SCALE behind its norm. The kernel strides over rows, so the grid can be capped at 65535 workgroups and a subgroup size other than the one the host assumed still covers every row. GGML_VK_NORM_SMALL=0 restores the old path (and disables both fusions); GGML_VK_FUSE_NORM_EXTRA=0 keeps the kernel but disables the two fusions. gfx1151 (RADV), same binary with GGML_VK_NORM_SMALL off/on, ABBA x2, -r 2 (mean of 8): Flash-Next REAP-320 pp2048 -ub 2048 773.3 -> 785.0 t/s (+1.5%), Qwen3.6-35B-A3B pp2048 -ub 2048 2100.4 -> 2145.5 (+2.1%), Qwen3.8-27B pp512 -ub 256 430.5 -> 431.6 (flat). Norm ops 640/640 in all three gate settings. PPL moves only by the reduction order (subgroupAdd vs the shared-memory tree). Only tested on gfx1151. Assisted-by: Claude Opus 5.5
updated readme.md to remove halo-box fork
vulkan : multi-row coopmat1 flash attention for prefill, with a V^T scratch
The cm1 int8 MMQ epilogue wrote its f32 result one element at a time. On the RDNA3/3.5 wave32 accumulator layout lane l holds column l % 16 and rows (l / 16) + 2e, so every store instruction scattered 4-byte writes over 16 columns. Two subgroupShuffleXor(16) per four elements now give each lane four adjacent rows of its column, and the lane writes them as one vec4. This covers MUL_MAT and MUL_MAT_ID. - The packed path needs stride_d (and batch_stride_d or the batch offset) to be a multiple of 4. Otherwise the per-element store runs, for example with odd M. A vec4 that crosses the row end falls back to scalar stores, so nothing is written past M. - It is enabled by spec constant 12 (PACKSTORE) of mul_mmq_cm1.comp. GGML_VK_MMQ_PACKSTORE=0 turns it off. RDNA4 (blocked layout) and other tile shapes keep the old store. - test-backend-ops: new int8 MMQ cases with ragged M/N (M = 20, 36, 132, 199, 68x2 batch; N = 67, 129) for MUL_MAT and MUL_MAT_ID. Results on gfx1151 (RADV): - MUL_MAT + MUL_MAT_ID 14980/14983 with the gate on and off; the 3 failures are the same pre-existing CONCAT k=5120 cases in both arms. - Positive controls: swapping lanes inside the vec4 fails thousands of MUL_MAT and MUL_MAT_ID cases with the gate on and none with it off. Dropping the row-end check fails the ragged-M cases, while odd M = 199 passes because it takes the scalar fallback. - llama-bench, same binary with the gate off/on, ABBA x2, -r 2 (mean of 8): Flash-Next REAP-320 pp2048 -ub 2048 769.5 -> 796.3 t/s (+3.5%), Qwen3.6-35B-A3B pp2048 -ub 2048 2100.5 -> 2187.1 (+4.1%), Qwen3.8-27B pp512 -ub 256 430.9 -> 441.5 (+2.4%). Every gate-on sample beat every gate-off sample. - Flash-Next perplexity is identical per chunk with the gate on and off. Assisted-by: Claude Opus 5.5
MUL_MAT_ID followed by MUL with the router weights ([1, n_expert_used, n_tokens]) was fused only on the mat-vec path. At prefill the expert outputs were written as f32, read back by a separate MUL, scaled and written again. The int8 MMQ coopmat1 MUL_MAT_ID pipelines get a sixth binding holding the weights. Bit 1 of the hoist_row_ids push constant tells the epilogue to scale each output row by its (expert slot, token) weight before the store. The product is the same f32 multiply the separate MUL does, so results are identical. The fusion is admitted only when ggml_vk_mul_mat_id_q_f16 will take the int8 MMQ path (checked against the pipeline's binding count); other MUL_MAT_ID GEMMs keep the separate MUL. GGML_VK_MMID_EPILOGUE=0 disables it. The packed vec4 store path scales its four rows by the same column weight. test-backend-ops: 9 new MUL_MAT_ID_FUSION cases (q8_0, iq4_nl, q4_K; n_used 10/4; ragged and odd M). MUL_MAT_ID + MUL_MAT_ID_FUSION + MOE_PREFILL 7561/7561 with the gate on and off. Scaling by twice the weight in the new epilogue fails every GEMM-path fusion case with the gate on and none with it off. gfx1151 (RADV), pp2048 -ub 2048, same binary with the gate off/on, ABBA x2, -r 2 (mean of 8): Flash-Next REAP-320 775.7 -> 800.2 t/s (+3.2%), Qwen3.6-35B-A3B 2100.4 -> 2207.5 (+5.1%). Every gate-on sample beat every gate-off sample. Perplexity is identical in every chunk. Assisted-by: Claude Opus 5.5
…th chunk The compute buffer is split into chunks of at most the backend's max size (1 GiB on Vulkan), but the chunk list was a fixed array of 16, and the 16th chunk was made unbounded so it could take whatever was left. A compute buffer past ~16 chunks therefore put everything remaining in one chunk that can be larger than the device allows: Flash-Next at -c 262144 on a 128 GB Strix Halo needed a 19.5 GB compute buffer and failed on a 4.4 GiB chunk, over RADV's 4 GiB buffer limit (issue 22). The chunk lists of the allocator and of the virtual buffer now grow (16, then doubling), and every chunk stays at the max size unless one tensor alone is larger. Graphs that fit in 16 chunks are laid out exactly as before. test_not_enough_chunks now also checks that every chunk stays within max_size; it fails on the old layout. Assisted-by: Claude Opus 5.5
vulkan : GDN conv straight from x and the state, and a short-row RMS norm with two fusions
…unks ggml-alloc : grow the chunk list on demand instead of an unbounded 16th chunk
vulkan : vec4 store for int8 MMQ on RDNA3, and MoE router weights in the MUL_MAT_ID epilogue
server : add --spec-draft-max-slots to stop drafting under concurrency
qsa : enable the incremental QSA cache on Vulkan
Only queries in the current stream consume the tail member table. Do not insert foreign image cells with repeated positions into that table. Retain shared cells visible to any query. Extend the existing QSA prefix test to exercise the production scan with foreign image cells, shared prefixes, multiple query sequences, and both compact and F32 metadata. Assisted-by: Codex
iTechMedic
added a commit
to iTechMedic/strix-llama.cpp
that referenced
this pull request
Oct 11, 2026
Preserve production candidate c2aae46 alongside PR halo-box#163. Production code merges unchanged. Resolve the test-file insertion conflict by retaining all PR halo-box#163 cases and adding the unchanged PR halo-box#165 regression and its main call. Combined CPU validation follows; full-model combination validation and daily promotion remain separate. Source-PR: halo-box#165 Assisted-by: Codex
3 of 5 tasks
This branch has not been deployed
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Overview
With several conversations in one unified KV cache (
--parallel N --kv-unified), qwen4exp can abort inllama_memory_hybrid_idx::set_input_qsa_scan:Reproduced trigger: one conversation holds an mrope image (a screenshot), and the full QSA scan builds the compact
(maskless) metadata for a text-only ubatch of another conversation. The scan runs:
LLAMA_QSA_INCREMENTAL=0;ubatch, or a sequence with shared (
seq_cp) cells.So the abort is also reachable with default incremental QSA when several conversations share a unified cache. Reproduced
case: one slot holds a screenshot conversation while two text conversations decode together (
--parallel 3 --kv-unified). The ordinary single-slot configuration is not affected by this trigger and was not seen to crash.Cause:
group_cellskeys a block on (sequence set, position block). An mrope image puts many cells of one sequence at oneposition, so they share one group and one slot.
group_members) walked every grouped cell, including other sequences' image cells, and asserted onecell per slot.
cells.seq_has(cell, seq_id),seq_id = ubatch->seq_id[i][0]). The foreign cells were never read; they only tripped the assertion.Change (6 lines):
the tail reader already use.
group_members.every skipped cell is one the reader rejects.
The existing
test-qsa-prefixgets one case. It runs the productionset_input_qsa_blocksscan on real cache cellmetadata (zero-layer model, no weights) with:
Unmodified master aborts at the assertion; this PR passes. The F32-bias-with-tails variant exercises
set_input_qsa_blocksdirectly; the qwen4exp graph uses tails only with the compact limits.Measurements
Baseline and candidate are separate builds from the same builder image and flags (ROCm) or the same host toolchain
(Vulkan). The baseline binaries were built the day before (2026-10-09) from the same image; both were run in this session.
Baseline / After (
llama-bench -p 2048 -d 0,12000,32000,64000 -b 4096 -ub 4096 -n 0 -r 4 -ngl 99 -fa on -ctk f16 -ctv f16 --load-mode none, order base, change, change, base; all samples kept):With
LLAMA_QSA_INCREMENTAL=0the changed loop runs for every QSA ubatch (same order):Production MTP server at 32K (shared Q8_0 head, draft n-max 3, vision, 1 slot, b/ub 4096,
--cache-ram 8192): llama-benchy 0.4.0--pp 2048 --tg 128 --depth 32000 --runs 5 --enable-prefix-caching, greedy, thinking off. Order: base, change, base,change, base, change; 3 servers per build.
numpy.random.seed(5151)) gave every serverthe same prompts. The prompt and request-payload hashes match on all 6 servers.
other by about 1.2 % on the same requests.
and draft contexts came from the incremental tracker.
server-to-server variation.
used different random prompts on each server, and acceptance depends on content; the seeded comparison above replaces it.
No speedup or memory change is claimed.
Correctness (target scope; this is not the full repository correctness flow, see the last section):
LLAMA_QSA_INCREMENTAL=0(ROCm: both orders; Vulkan: image first)--parallel 3 --kv-unified(ROCm: MTP and serial; Vulkan: MTP)n=2(shared prompt cells) next to a screenshot conversation,LLAMA_QSA_INCREMENTAL=0(ROCm)A diagnostic-only build (not part of this PR) also checked the compact QSA metadata against an independent brute-force
reference of the selection policy: about 4,600 metadata sets, including foreign screenshots, multi-row ubatches and
shared cells, with 0 mismatches. This validates the QSA block/tail metadata. It says nothing about the quality of
generated answers.
Additional information
307c50d26). The tail-fill code is unchanged from that revisionthrough b1015 to current master, and the reproduced foreign-image trigger above still crashes current master. Eval bug: Qwen4exp crashes when -np > 1 and -kvu is set. PR #86 Regression. #125's
original trigger was not reported in enough detail to confirm it is this one. This PR removes the crash for foreign
sequences' cells; it does not claim to fix every possible cause of this assertion. A duplicate slot among cells that a
query can read would still assert.
scan for many ubatches but does not change it.
was reproduced on a qsa : keep the incremental QSA prefix across mrope images (target and MTP draft) #163 build with the incremental cache off. This PR is independent of qsa : keep the incremental QSA prefix across mrope images (target and MTP draft) #163 and applies to master
alone.
Requirements
this fork.
Assisted-by: Codex).description.
src/llama-memory-hybrid-idx.cppandtests/test-qsa-prefix.cpp. GLOBAL PASS is not claimed.n=2/ long-sequence server runs, no Vulkan performance runs.sequence that shares a ubatch with it. On ROCm, byte identity is therefore shown for text sequences and controls
only.
response validity and the metadata check, not by token identity.
--cache-dirwas not exercised. Partialseq_rmis refused by this hybrid model,so rollback was exercised through state save/restore, MTP draft rejection and RAM-cache restores.
🤖 Generated with Claude Code
https://claude.ai/code/session_01FqHrSzr2sAmQUm8syTy1ES