Fix typo in README by marckohlbrugge · Pull Request #45 · ggml-org/llama.cpp

marckohlbrugge · 2023-03-12T11:10:37Z

No description provided.

QVAC-6569: Redirected performance logging to GGML_LOG

Temp fix for multithreading bug

* add flash attn op * expend src tensor size * add flash attn sources * add quantize row functions * make a separated file for vec_dot * wip * wip * refactor: rename quants.hpp includes and add vec_dot to type traits * add flash_attn impl * split vec_scale_f32 * move vec_reduction_qf32 to vec_ops * add vec_scale_f16 * opt * add vec_mad * implement vec_mad_f16 * opt * add op template * opt * add align version * enable flash attn * wip * log print improve * add profiler log * wip * wip * add multi sub proc perf tracker * increase log buffer * remove sub prov pcycle * wip * wip * add prefetch for vec_dot * wip * wip * opt f16 vec dot * opt f16 vecdot * reuse vec_dot_product_impl in vec dot f32 * small opt to unblock pipeline * opt on aligned address wip * Revert "opt on aligned address" This reverts commit 27be1eb. * add profiler log at thread_pool * wip * invalidate all... * Reapply "opt on aligned address" This reverts commit f075a4c. * add is_constant for tensor config * disable align tensor opt in mul_mat * wip * wip * vec_scale_impl: unrolling the loop * wip * wip * replace reinterpret_cast with direct pointer access for write/read buffers * add fetch * wip * wip * wip * add log * check tensor shape at flash_attn * wip * wip * fix: update tensor type handling in flash_attn_impl * wip * fix: align cache size * fix: qf16->hf * fix: swap order of elements in vector combine for correct scaling * fix: opt f16 scale and mad * fix leftover fetch * wip * load into vector pair * opt cache size calculation in flash_attn_impl * refactoring: hold vtcm at thread local object * wip * add profiler log * mark tensors as modified * restrict tensor invalidation to the first thread in compute_impl * Revert "restrict tensor invalidation to the first thread in compute_impl" This reverts commit 0a8ff2b. * invalidate last tensor in compute_impl * invalidate last tensor in compute function * wip * refactor dequantize_row_q4_0 to simplify vector alignment * wip * refactoring: move VTCM quota calculation to thread pool * wip * fix: correct condition check for HEXAGON_SDK_ROOT existence * wip * wip * wip * wip * fix: update condition checks match the naming * fix: improve tensor handling checks and logging in graph and operation implementations * wip

- turbo4 K+V results on Qwen3.5-27B (-0.32% vs q8_0) and Qwen3-14B (+6.3%) - Sparse V dequant benchmarks: MoE native dequant +10.9% at 8K - Gemma-3 turbo3 results post-iSWA fix (+3.3%) - KVLinC no-K-rotation negative result - Speculative decoding negative result - CUDA 13.2 compatibility verified - Experiments #31, #39, #42, #45, #49, #50, #51 status updates Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.com>

turbo4 SET_ROWS was using turbo3's shared template with wrong 2+1 bit packing. New dedicated kernel_set_rows_turbo4 with correct 3-bit packed indices + QJL signs. PPL: 679 → 6.19. Also added turbo4 prefill FA kernel instantiations (non-vec path). QJL ablation finding: disabling QJL improves PPL from 6.1894 to 6.1756 (identical to turbo3). QJL correction hurts quality in attention context. Consistent with scos-lab issue #45. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

turbo4 SET_ROWS was using turbo3's shared template with wrong 2+1 bit packing. New dedicated kernel_set_rows_turbo4 with correct 3-bit packed indices + QJL signs. PPL: 679 → 6.19. Also added turbo4 prefill FA kernel instantiations (non-vec path). QJL ablation finding: disabling QJL improves PPL from 6.1894 to 6.1756 (identical to turbo3). QJL correction hurts quality in attention context. Consistent with scos-lab issue ggml-org#45. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

* iq1_tn: adding CUDA dequantize * iq1_tn: adding CUDA dot product * Delete commented out stuff * Delete forgotten TODO --------- Co-authored-by: Iwan Kawrakow <iwan.kawrakow@gmail.com>

turbo4 SET_ROWS was using turbo3's shared template with wrong 2+1 bit packing. New dedicated kernel_set_rows_turbo4 with correct 3-bit packed indices + QJL signs. PPL: 679 → 6.19. Also added turbo4 prefill FA kernel instantiations (non-vec path). QJL ablation finding: disabling QJL improves PPL from 6.1894 to 6.1756 (identical to turbo3). QJL correction hurts quality in attention context. Consistent with scos-lab issue ggml-org#45. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

…ss + hardening) Hardening follow-up to ggml-org#39 — same Hopper sm_90a wgmma path. Adds: device-keyed repack cache + mutex + publish-after-sync, per-device attr_set/NSM (multi-GPU correctness), CUDA_CHECK on allocations, defensive int8 clamp, target-scoped CMake. Verified: perf-neutral vs ggml-org#39 (worst -0.5%), +7.6-8.9% vs stock MMQ holds, 86/86 test-backend-ops; multi-arch/multi-GPU sweep green (2xH100 dual-GPU, L40S no-regression, H200 +6%). Closes the multi-GPU gap left open in ggml-org#39.

Support VoxCPM 0.5B, 1.5, 2 inference on llama.cpp

@briansp2020

Campaign fold (wip/moe-expert-cache, beta2): adds the generic backend expert-cache interface -- `moe_cache_update`, `moe_cache_take_over`, `moe_cache_promote` and `moe_cache_gather` in the backend iface (ggml-backend-impl.h) -- so a backend can drive a residency policy for a host-resident routed-expert table, plus the scheduler half of the consumer: the routed-expert rebalance onto the layer's owning device and the merged per-layer MoE split (ggml-backend.cpp), the device-side expert gather that defers a gatherable routed table from whole-tensor H2D staging, the input-takeover fast path with its deferred promotion pass, the Meta delegation to the per-device simple backends (ggml-backend-meta.cpp), and the `#define GGML_ENV_STR` cache macro the campaign's gates use. The cache engine and its CUDA consumers land in block 13; the gemma4 `-sm tensor` exclusion is in block 14. The CPU and RPC backend vtables name the four new iface fields (NULL) so the set stays warning-free. Amended for the promotion candidate (beta4): keep the scheduler changes prefill-neutral. The device gather now serves only the below-`sched_stage_min_tokens` band (above it the r16/r26 whole-shard staging wins: Q4_K_M ub8192 staging 5283 vs gather 4035; `-sm layer` ncmoe 40 4511 vs 2625), and the routed-expert rebalance is gated to the decode/verify band (`MUL_MAT_ID` `ne[2] <= 8`), where the cache runs. Prefill then matches `main` (r28) at every measured cell while the below-gate gather wins (ub512 +17 %, ub1024 +19 %) are kept. beta5 (2026-10-02): the device gather now also serves a large (>= 224 MiB) host-resident expert table above the width gate. Re-measured with --lazy-mode off (the qwen4exp PLE mmap warm-up confounds every prefill comparison): qwen4exp 450 MiB ub8192 gather 3052 vs staging 1403; Q8_0 272 MiB 3595 vs 3552; Q4_K_M 144 MiB gather 4291 vs staging 5624, so Q4_K_M keeps the width gate. r35 (2026-10-04): issue ggml-org#87's synchronous graph-input fix is now default-on. The r33 `GGML_SCHED_SYNC_GRAPH_INPUTS=1` A/B candidate still reads the variable, but unset means enabled; `=0` restores the r26 async behaviour. The reporter confirmed it clears their recurrent-state H2D race, and same-seed output is unchanged both ways. r5 (2026-10-05, PR ggml-org#96): re-stage a host-resident expert table for a second MUL_MAT_ID consumer in a later split. `ggml_backend_sched_split_graph` registered a weight as a split input only when it first created the copy, so when the same expert weights fed `MUL_MAT_ID` in two splits the later split reused the earlier split's copy. With the expert cache armed that copy was taken over by the 1-row decode-band consumer (`moe_cache_take_over` aliases it and the op never fills it), so the later wide consumer read stale bytes. qwen4exp's unmasked MTP export (a full-row last-layer FFN next to the gathered logits tail) is exactly that second consumer: the wide op produced NaN in `t_h_nextn`, which reached the drafter's KV and collapsed MTP draft acceptance across every slot. When the copy already exists and the node is a `MUL_MAT_ID` reading WEIGHTS through src[0], register the weights as an input of the current split as well (once per split), so they are staged for this split's routing. Single-consumer graphs are unchanged. Contributor PR ggml-org#96 (@briansp2020). Folded 2026-10-07 (TODO ggml-org#45): the stale `--load-mode none` + `-sm tensor` host-expert warning is removed. Its stated mechanism (host experts landing in a pageable CPU_REPACK buffer) no longer holds: the loader pins the routed experts in `ROCm_Host` and the fault no longer reproduces (14/14 clean on r24). Record: archive/work/host-pinned-buffer-crash/.

@briansp2020

The RDNA memory campaign, squashed into one block: six validated wins that remove ~3.4 GiB/GPU of compute-buffer VRAM and ~1.2 GiB of host buffer on qwen4exp (ctx 204800 / ub 2048 / q8_0 KV), plus ~800 MiB/GPU + 800 MiB host on every dense model, at a measured ~1.3 % prefill / ~0.3 % decode cost. Each win keeps an environment kill-switch for A/B (the V4 arm is opt-in). W1 - QSA score-chain memory (GGML_QSA_SCORE_MEM, default 1) relu before the 4-D reshape (elementwise, bit-identical, but it stops the allocator from holding the mul_mat result and its relu simultaneously) + n_blocks-chunked assembly with ggml_concat past a size threshold. qwen4exp ub 2048: 6690.40 -> 4450.40 MiB/GPU. W2 - derived QSA per-block bias + derived visibility (GGML_QSA_DERIVED_BIAS, GGML_QSA_DERIVED_VIS, LLAMA_QSA_SPARSE_FA, defaults 1/1/sparse): the n_blocks x n_tps F32 bias tensor and the additive kq mask are no longer materialised; the fused top-k/FA kernels derive both in-kernel from compact per-cell state. Includes the input-fill null guards (a tensor the graph never consumes is left unallocated by the gallocr, so every fill site must tolerate buffer == nullptr) and the llm_graph_input_attn_k null-mask guard. qwen4exp ub 2048: 4450.40 -> 3251.39 MiB/GPU, host 1262.70 -> 63.69 MiB. W3 - keys-only QSA indexer cache (LLAMA_QSA_KEYS_ONLY, default 1): the indexer scores blocks of keys and never reads a stored value, so its cache is created with no V tensor. Indexer KV 956.26 -> 318.76 MiB/GPU at ctx 204800 (the store is replicated on every GPU). W4 - ggml-alloc: release view sources whose views are never consumed (the counting pass increments view_src->n_views for every view node, but the free pass only decrements it when the view is released - a view with no consumers is never released, so the inflated count blocks the view source's release and in-place reuse). Repro 56.00 -> 16.00 MiB. No gate: it is a bug fix, A/B with beta/block-15-campaign-wins/ab/w4-revert.patch. V3 - derived kq mask for the plain attention path (LLAMA_KQ_MASK_DERIVED, default 1): the packed n_kv x n_tps F16 mask and its host mirror are no longer materialised - the MMA FA kernel derives each cell's visibility from compact per-cell state (GGML_OP_FLASH_ATTN_EXT gains src[5..7]: cell positions, per-token hi/lo bounds). The packed tensor is still created in every graph, so it simply ends up with no consumer: the allocator leaves it unallocated and the existing fill guards skip the fill (no model list, no consumer can be mis-served). The predicate is proven cell-by-cell value-identical (including both SWA cache variants and non-causal) and gates itself off on every unsupported case (alibi, multi-sequence, small/verify batches, non-MMA backends, M-RoPE ext clause). Dense models: -799.20 MiB/GPU and -799.21 MiB host; gemma-4-E4B (ISWA) -809.18/-809.18; gemma-4-31B (ISWA) -811.17/-811.18; scaling exactly n_kv x n_tps x 2 B. The backend support probe is skipped entirely on a multi-stream KV cache (n_seq_max > 1 without kv_unified, and deepseek4, which keeps per-sequence streams even in unified mode): it forces a single-sequence probe graph, and on such a cache deepseek4's indexer stream layout asserts, while the derived form is unreachable there anyway (kq_mask_derivable() requires a single stream). V4 - native q8_0 K/V in the flash-attention kernels (GGML_CUDA_FA_KV_NATIVE, OPT-IN, default 0): the whole-cache F16 staging scratch and its per-ubatch conversion are replaced by dequantising each 16-byte staged chunk while the shared K/V tiles are loaded (a chunk is 8 elements = a quarter q8_0 block, and every MMA/TILE batch is a multiple of 8 elements, so a chunk never straddles a block). The staged values are bit-identical to dequantize_block_q8_0_f16. Dense models: -744 MiB/GPU (4B), -632 (27B), -1224 (gemma-4-31B) at ctx 204800/ub 2048, more at smaller ub; qwen4exp unchanged. Ships opt-in: it costs ~1.7-1.9 % prefill (the lost cp_async pipeline - a quantized source cannot be copied asynchronously) with a large memory win and no cheap way to close the gap; decode is untouched. TILE and VEC already read bf16/f16 natively. Validation (2026-09-10, 3x R9700/RDNA4, ctx 204800, q8_0 KV unless noted; full record: beta/block-15-campaign-wins/): - reserve matrix, all five models x ub 2048/1024/512 x V4 off/on: every number matches the per-win records (the wins compose additively: -2240 W1, -1199 W2, 0/-637.5 W3, -799 V3, -744/-632 V4 per GPU). - same-seed generated text byte-identical on 4B, 27B, gemma-4-E4B (ISWA), gemma-4-31B (ISWA) and qwen4exp across every gate combination (V3 on/off x V4 on/off; qwen4exp W1/W2/W3/W4/V3/V4 flips - 7 configurations), at a short and a 40k-token prompt. - adaptive-MTP gate unchanged: 27B inline draft 0.76744 (66/86, mean 3.28) identical in all four gate combinations; qwen4exp draft 0.44262 (54/122) identical in all six, and equal to the block-14 baseline; MTP still +26 % over plain decode. - test-backend-ops: FLASH_ATTN_EXT on ROCm0 (both V4 gates) and CPU, the derived FA cases, VIEW/CONT/CPY/DUP/CONCAT; test-alloc and test-batch-alloc; the W4 repro goes 56.00 -> 16.00 MiB and the revert patch restores ggml-alloc.c byte-identically. - prefill cost, interleaved same-binary A/B: V3 -1.28 % pp20480/ub 2048 on the 4B (the 27B is neutral, +0.28 %), V4 a further -1.85 % (4B) / -1.72 % (27B); decode within noise in every case. Known pre-existing issue found during this validation (NOT introduced here, reproduces on block 14): gemma-4-E4B-it on 3 GPUs with -sm tensor aborts in the meta splitter (ggml-backend-meta.cpp:1177, a FLASH_ATTN_EXT node whose K source has zero extent on one buffer because n_head_kv = 2 < 3 devices). It works on 1 GPU, on 2 GPUs, and on 3 GPUs with -sm layer. All other models are unaffected. Out of scope for this block. Beta staging, gate table and the full validation record: beta/block-15-campaign-wins/README.md; the per-win designs: wip/arch-independent-memory/{V3-DERIVED-KQ-MASK-PLAN,V4-NATIVE-Q8-KV-PLAN, DERIVED-MASK-DESIGN}.md and wip/qwen4exp/qsa-memory/. Amended 2026-09-10 with V5 - native bf16 K/V in the MMA flash-attention kernel (gated by the SAME GGML_CUDA_FA_KV_NATIVE switch as V4, default 0, i.e. opt-in): a bf16 K/V cache no longer needs the F16 staging scratch. A bf16 source row and the F16 tile it feeds have the same byte layout (2 bytes per element, 16-byte chunks), so the MMA loader converts each staged chunk in registers (bf16 -> f32 -> f16, bit-identical to the launcher's ggml_get_to_fp16_cuda(GGML_TYPE_BF16)) instead of copying it, and the scratch sizing plus the whole-cache conversion pass are skipped for that operand. The staging source is now a single type code (fattn_kv_native_t{FATTN_KV_NATIVE_NONE,Q8_0,BF16}) shared by the launcher, the alloc-size query and both kernels, so they cannot disagree on whether the scratch exists (the old use_q8_K/use_q8_V flags are subsumed). - reserve (ctx 204800, ub 2048, bf16 KV, arm on): 4B 968.86 -> 256.86, 27B 1072.86 -> 488.86, gemma-4-E4B (ISWA) 1062.89 -> 404.89, gemma-4-31B (ISWA) 2068.89 -> 716.89 MiB/GPU - exactly the F16 cache's footprint; ub 1024/512 win 756/778 (4B) and 746 (27B); qwen4exp is unaffected (its FA path never staged bf16 - f16 == bf16 == on/off there), and the TILE/VEC paths already read bf16 natively (ub 8: 8.09 both ways). - same-seed generated text byte-identical (bf16 native vs staged vs f16) on 4B, 27B, gemma-4-E4B and gemma-4-31B, short and 3k/40k-token prompts, with the V3 derived mask active; the two wins are additive (4B: derived mask -799.2 MiB, bf16 scratch -712.0 MiB, both verified independently). - adaptive-MTP gate unchanged: 27B 0.82716 and qwen4exp 0.44262 identical with the arm on and off; the recorded q8_0 references (27B 0.76744, qwen4exp 0.44262) are unchanged by the refactor. - test-backend-ops FLASH_ATTN_EXT: 7859/7859 on ROCm0 and CPU; the 2704 bf16 K/V cases (every head size, including the 576/512 MLA V-is-K-view layout) and the 365 q8_0 cases pass with the arm on and off, with identical case lists. - prefill cost (interleaved same-binary A/B, arm off -> on): the conversion itself is free (native bf16 staging measures within 0.2 % of an F16 cache), but dropping the scratch costs 4B -0.2 % (pp2048), +0.3 % (8192), -1.06 % (20480), -2.36 % (40960) and 27B -0.76 % (20480); decode is within 0.1 % everywhere. The launcher's F16 copy is a dense, normalized copy of the cache view (the GQA heads are interleaved - for the 4B nb[1] is 2048 B for a 512 B row), and the native path re-reads that view on every staging pass. Hence opt-in, exactly like V4 (D9: a sub-2 % loss with a large memory win and no cheap fix ships behind a switch). Mixed K/V cache types (bf16+q8_0, f16+q8_0, either direction) remain a pre-existing trap, documented but not fixed here: the attention falls off the GPU kernel path (graph splits 18 vs 2, ~1.5 GiB host compute buffer, 4B pp2048 7924 -> 640-1049 t/s). Same-type K/V is the practical choice. 2026-09-10 RDNA3_5 (gfx1151) amendment: V3 (derived kq mask) now engages on a ROCm/HIP iGPU. The probe rejected GGML_BACKEND_DEVICE_TYPE_IGPU as "not the CUDA/HIP backend", so on Strix Halo V3 was silently disabled and its ~800 MiB compute + ~800 MiB host win was lost. The device-type gate now accepts IGPU (the ROCm/HIP reg name is still required), and the derived form additionally requires a single KV stream: build_attn_mha derives the attention stream count from the KV tensor (k->ne[3] == the cache's n_stream), so with n_seq_max > 1 -- even for a single-sequence ubatch -- the derived tok_lo was built for one stream while Q was reshaped for n_seq_max streams and the shape assert in ggml_flash_attn_ext_add_kq_derived aborted context creation. kq_mask_derivable() now rejects n_stream != 1, so a multi-slot context keeps the packed mask (no crash) and a single-stream context gets the win. Validated on Strix Halo (gfx1151, ROCm 7.14): V3 on/off byte-identical over 7 KV/arm configs x 16 runs x 129 tokens (2064 gate cells each); reserves reproduce the RDNA4 numbers exactly (4B ctx 204800/ub 2048: V3 -799.20 MiB compute / -799.21 host, V5 bf16 968.86 -> 256.86, V4 q8_0 1001.13 -> 257.13); the derived test-backend-ops cases pass on ROCm0; and llama-server --parallel 4 (the config that aborted before) now serves and passes the 16-run gate. Discrete GPUs are unaffected (they report GGML_BACKEND_DEVICE_TYPE_GPU and normally have n_stream == 1). Amended 2026-09-11 (9th re-cut): fix a variable-shadowing bug in build_attn_qsa that broke the LLAMA_QSA_SPARSE_FA=0 dense arm for every model/type. The V2/V3 refactor wrapped the top-k mask chain in `if (kq_mask != nullptr) { ... }` and added an outer `kq_mask_top_k` declaration, leaving the chain's own `ggml_tensor * kq_mask_top_k = ggml_set_rows(...)` inside the block as a NEW local. The chain was therefore built (whenever the mask existed) but its result never reached the attention: build_attn_mha got the outer `nullptr`, the chain's nodes were unreachable from the graph output (so ggml never emitted them), the packed mask lost its only consumer (so the allocator dropped it and set_input_kq_mask was skipped), and the dense arm attended with no mask at all - a full causal leak. Observed as qwen4exp perplexity ~1.02-1.06 on every KV type (delivery: 6.49-6.55), and as ~1.03 on RANDOM text (delivery ~18.4), with dense-arm greedy text differing from the delivery. Fixed by assigning the block's result to the outer variable (dropping the inner `ggml_tensor *`). The production (sparse) arm never enters that block, so it is bit-identical before and after. See beta/block-15-campaign-wins/BETA-TESTING.md 4c and WORKLOG.md 2026-09-11 (11). 2026-09-14 amendment (issue ggml-org#30, "wider configuration"): V4 is no longer opt-in for the sub-F16 quantized cache types, and q4_0 gains a native arm. - GGML_CUDA_FA_KV_NATIVE is now a three-state policy: unset = auto (native q8_0/q4_0 ON, native bf16 OFF), =1 forces both on, =0 forces the pre-amendment F16-staging path (the escape hatch). The whole-cache F16 staging pass is proportional to n_kv and runs on every decode step, so removing it for the sub-F16 quants is worth +23 % tg64 at d65536 on gfx1201 (q8_0 18.92 -> 23.29 vs stock 22.43) at ~1.2 % prefill, and it recovers ~744 MiB/GPU of scratch. - Native q4_0 arm beside q8_0/bf16 (ggml_cuda_fattn_dequantize_q4_0_chunk, arithmetic-identical to convert.cu's dequantize_block_q4_0) in the tile and MMA kernels: q4_0 d65536 19.72 -> 22.82 vs stock 21.03, ~1.3 % prefill. - Bit-identical to the staging path (same-seed greedy text) and width-pure (W = 1..8 one logits hash per KV type); MTP n_max 3 acceptance unchanged (0.75182). Full evidence: wip/issue-30-mtp-decode-regression/MEASUREMENTS.md. - Side effect: also fixes the adaptive-MTP high-context load failure (--spec-draft-n-max 12 -c 196608 q8_0, default 4 slots): the F16 staging scratch was exactly the 260 MiB the draft context was short, so it now loads and generates. 2026-09-15 amendment (issue ggml-org#30, the reporter's r3 review): the q4_0 NaN fix, the prefill band split, and native arms for the last four quantized K/V types. - Mixed-K/V kernel contract (the reporter's 4 NaN failures in test-backend-ops -o FLASH_ATTN_EXT). The tile kernel is instantiated with ONE type_KV for both operands and builds its native descriptors from that compile-time type, while launch_fattn chose its native read per tensor; a mixed pair (K=q4_0, V=f16) therefore fell back to the F16 tile with K's staging skipped, and the kernel read raw q4_0 as F16. launch_fattn now takes the kernel's native type explicitly (kv_native_kernel; the tile passes its type_KV, the vec passes NONE, the MMA keeps per-operand semantics). llama.cpp rejects mixed caches, so only the op test could see it. - ggml_cuda_flash_attn_ext_get_alloc_size's TILE case never learned the q4_0 arm, so a q4_0 cache kept reserving the F16 scratch the launcher no longer used: -c 196608 q4_0 849.04 -> 123.04 MiB -- the arm's memory win had never actually been delivered. - Prefill band split: a prefill (n_q > 8) stages the operand into F16 while decode/verify (n_q <= 8) keeps the native read, recovering the quantized-KV prefill cost (q8_0 pp150000 661.0 -> 691.4 on 1 GPU, 996.0 -> 1076.9 on 2, 1111.4 -> 1199.0 on 3; q4_0 694.5). The staging scratch comes from a new per-context, per-stream arena (ggml_backend_cuda_context::fattn_stage / fattn_stage_get) instead of the compute-graph reserve, which sized it for n_ctx (~726 MiB at a 200k context) even though the real scratch tracks the prefix; a multi-token graph is never CUDA-graph captured, so a runtime allocation there is safe. Bounded by GGML_CUDA_FA_STAGE_MAX_MB (default 512 MiB, 0 = unbounded). Arch-gated: RDNA3_5 keeps the native read at prefill (it wins at every depth there: gfx1151 9B q8_0 pp16k 1410 vs 1405, pp65k 1059 vs 1043), RDNA4/RDNA3_0 stage. - Native arms for q4_1 / q5_0 / q5_1 / iq4_nl (ggml_cuda_fattn_dequantize_{q4_1,q5_0,q5_1, iq4_nl}_chunk, arithmetic-identical to convert.cu), closing the last gap in the V4 set -- the four types stage through F16 otherwise. tg64 @ d32768 staged -> native: gfx1201 +9.9/+10.8/+12.5/+8.8 % (q4_1 23.14 -> 25.44, q5_0 22.16 -> 24.56, q5_1 22.23 -> 25.00, iq4_nl 22.92 -> 24.94) with prefill unchanged; gfx1151 (9B) +22.9/+26.0/+26.7/+22.0 % for a 0.6-1.1 % prefill cost. Two bugs the op test caught on the way: the tile loader chose the native branch with a hand-written `type_KV == Q8_0 || Q4_0` test (the new instantiations then took the F16 branch and read an unwritten staging buffer -- the hsk=72 NaNs), and the q5_0/q5_1 chunk helpers took the low nibble in both halves. - Same-seed greedy text is bit-identical native vs staging for all eight KV types on both arches, and test-backend-ops -o FLASH_ATTN_EXT is 5951/5951 on both. Width purity: gfx1151 pure for all eight types; gfx1201 pure except the documented pre-existing q4_0 logits-level band edges (GREEDY-PURITY.md 36, whose guarantee is now stated at the text/acceptance level). 2026-09-15 amendment (build time): the tile kernel's native-KV type axis is instantiated in the generated template-instance TUs again instead of implicitly in the tile dispatch TU. Block 03 turned type_KV into a template parameter of ggml_cuda_flash_attn_ext_tile_case so the tile kernel could read BF16 (and later the quantized types) natively, while DECL_FATTN_TILE_CASE / EXTERN_DECL_FATTN_TILE_CASES kept covering only F16/BF16. Because the dispatch in fattn-tile.cu has an unconditional `case` per native type, the other six types were then instantiated *implicitly in that TU*: 12 head-size combinations x 6 types = 72 case instantiations (96 symbols including the 24 externed F16/BF16), each pulling in both softcap variants and the whole ncols2 chain. With block 08's FA_QUANTS default (8 diagonal pairs) that single TU took 509 s of a 538 s `-j16` backend build on a 16-core machine (measured 2026-09-15, gfx1201) -- it *was* the critical path, and no amount of build parallelism could help. The macros now cover every type the dispatch can select, so the 12 generated template-instances/fattn-tile-instance-*.cu files instantiate 8 cases each and the dispatch TU holds only `extern` declarations. Build time: `ggml-hip` 538 s -> 330 s, and the tile dispatch TU itself drops from 509 s to under 10 s. The remaining critical path is the `fattn-mma-f16` instance set, and that one is ours too: this block's native-KV arms instantiate the whole WMMA kernel once per KV type *inside every instance TU* (the arm chain in `ggml_cuda_flash_attn_ext_mma_f16_case`), so each of the 21 instance files compiles 8 kernel copies. Measured on the same TU (`ncols1_4-ncols2_4`): object 0.90 MB -> 7.26 MB, compile 6.7 s -> 229 s. Fixing it needs either a finer generated-file granularity (~21 -> ~150 TUs) or a runtime KV-type dispatch in the WMMA loader, i.e. a code-path change with its own A/B, so it is left as a follow-up rather than folded in here. Nothing about the kernels changed -- same template arguments, same flags, same device code; only the TU that emits them differs. Verified: `test-backend-ops -o FLASH_ATTN_EXT` 5951/5951 with zero failures, the 27B prose text hashes bit-identical to the pre-amendment build (q8_0 472b282950b5, q4_0 118eb7f5fe85, f16 70960317a203), and `tg64@32768`/`pp8192` reproduce the recorded values within run-to-run noise across five KV types. Amended 2026-09-17: the HIP build now passes `-Wno-pass-failed`. The FA kernels carry bare `#pragma unroll` hints on runtime-bounded loops, so AMDGPU clang reports every one it cannot honour; the native-KV arms instantiate the WMMA kernel per KV type per instance TU and multiply that into ~10k "loop not unrolled" warning lines per clean build. The hints are advisory and the pass already failed, so the build output is unchanged -- only the flood is gone. HIP-only (the diagnostic is Clang/AMDGPU). See wip/build-time-regression/. Amended 2026-09-18: the tile template instances are generated per (head size, KV type) instead of per head size, matching the per-type axis this block introduced. Build-time only: the same 96 instantiations, identical symbols. Amended 2026-09-25 (issue ggml-org#45): the RDNA4 GQA-6 decode/verify flash-attention band. At head 256 with GQA 6, the tile kernel can only fold ncols2 = 2 (its ncols2 must divide the GQA ratio), so the whole n_q <= 8 band fetched and dequantized every K/V element once per head pair -- three times per query row at 276 GB/s on q8_0. The band now routes the whole band (n_q = 1 included) to the WMMA kernel with the GQA group folded into ncols2 = 8 (existing (256, ncols1 4|2, ncols2 8) instances, no new instantiations): one block covers the group, and the KV is split round-robin over a fixed P = nsm blocks per output tile (launch_fattn, GGML_HIP_FA_BAND_WMMA_SPLIT override). P depends on neither n_q nor the KV length, so decode and every verify width accumulate the same partials in the same order -- the round-robin form is what keeps the GREEDY-PURITY band invariant (a contiguous stream-k split moved its split points with the 256-padded KV length and broke it). - kernel chooser, ncols dispatcher and launch_fattn share one self-contained predicate (ggml_cuda_fattn_band_wmma_applies), and launch_fattn also requires the ncols2 == 8 template parameter, so a disagreement can never launch the 2-D grid into a kernel whose fast path is compiled out. - coverage follows ggml_cuda_fattn_kv_native_type, not a hand-written type list: every native quantized K/V type wins (q8_0, q4_0, q4_1, q5_0, q5_1, iq4_nl); f16 stays on tile (DRAM-bound) and bf16 stays on tile while its native arm is opt-in. - default ON (maintainer policy); GGML_HIP_FA_BAND_WMMA=0 opts out (=2 picks ncols1 = 2). GGML_CUDA_FA_WMMA_256=0 / GGML_CUDA_FA_WMMA_MAX_HEAD<256 still disable the band. Measured on gfx1201 (R9700), 27B qwen35 head 256 GQA 6, q8_0 KV, op level kv 16384 (tile -> band, us/run): n_q 3 326 -> 150, n_q 5 524 -> 225, n_q 8 814 -> 230; n_q 1 141 -> 148 (q4_0/q4_1/iq4_nl are faster at n_q 1 too). End to end, 27B UD-Q4_K_XL draft-mtp n3 at ~40k ctx, band off -> on: 1 GPU 43.3 -> 51.3, 2 GPU tensor 67.9 -> 76.6, 3 GPU tensor 77.6 -> 88.8 t/s (+13-18 %), while pp512/2048/4096 are flat (the band is gated to n_q <= 8, so prefill is untouched by construction). Purity re-verified with the band on: test-backend-ops -o FLASH_ATTN_EXT 6340/6340 on ROCm0 (389 GQA-6 cases across all eight KV types, both layouts); plain == draft-mtp byte-identical on all eight KV types at ~40k ctx, 3-GPU tensor (acceptance 0.83-0.86), and q8_0/q4_0 text identical across --spec-draft-n-max 3/5/7 (verify widths 4/6/8); 1 GPU layer, 2/3 GPU tensor and 3 GPU layer are all pure. New qwen35-band eval and perf cases added to tests/test-backend-ops.cpp. Folds beta/mmb-general: 0002 (qsa3 packed WMMA), 0003 (qwen4exp always-QSA flip), 0004 (HC16 native-BF16 producers), 0005 (fused indexer top-k), 0008 (HC16/beneficial defaults on ggml-cuda.cu + llama-context.cpp), 0011 (qsa3 on RDNA3_0), 0013 (hc_gate_mix), 0016 (-lzm auto + managed PLE gate), 0017 (HC BF16 streams), 0018 (out_xn BF16), 0019 (indexer relu+head-sum), 0020 (F32 activations on eval callback), 0021 (sparse MTP-draft attention), 0022 (derived indexer cache ON), 0023 (QSA gfx1151 crossover), 0024 (HC16 per-backend state), 0026 (meta child graph_optimize markings). Amended 2026-09-27 (issue ggml-org#53): the derived kq-mask inputs must not be read on the host, and the frontend test cases must set tensors through the backend API. The r10 fully-masked KV-group skip built its batch-wide bitmap by dereferencing `tok_lo`/`tok_hi` on the CPU in `launch_fattn`. Those are host graph inputs (GGML_TENSOR_FLAG_INPUT), and the backend scheduler copies such an input to the compute backend, so the tensors the launcher sees are the *device* copies. A host read of device memory is an access violation wherever the allocation is not CPU-mapped - on Windows/WDDM it is a 0xC0000005 in ggml-hip.dll on the very first prefill ubatch (the reporter's r10 crash) - and a race against the in-flight copy everywhere else. The batch-wide [lo_min, hi_max] window is now computed inside `flash_attn_kq_derived_blocks` from the device-resident `tok_lo`/`tok_hi`; the bitmap it produces is unchanged (same integer min/max), so the skip is still exact. Verified on gfx1201: test-backend-ops -o FLASH_ATTN_EXT 6354/6354 with the MMA and the tile kernel (and the 20 derived/mask_hole cases specifically), FLASH_ATTN_QSA 26/26, and 4B Q8_0 same-seed text identical across {skip on, skip=0, derived=0}. The same class of bug is fixed in tests/test-backend-ops.cpp: the derived (`cell_pos`/`tok_lo`/`tok_hi`), packed `mask_hole` and `FLASH_ATTN_QSA` (`idx`/`m`) initializers wrote `t->data` directly instead of going through `ggml_backend_tensor_set`, so on Windows the first `derived=1` case segfaulted before any derived case ran (also reported in issue ggml-org#53). The QSA tensors are block-14 test code; the fix rides in block 15 because it is the last block. Amended 2026-09-29 (issue ggml-org#65): cache the getenv() lookups on the fusion and staging hot paths. ggml_can_fuse_subgraph_ext() read LLAMA_HC_CN_DEBUG for every candidate fusion window (millions of calls per pass), ggml_cuda_try_fuse() read it again in the four hc_combine_norm matcher paths, and gdn_conv_enabled()/ple_conv_enabled() read GGML_CUDA_DISABLE_CONV_FUSION once per conv launch. On Linux getenv() is cheap, but on Windows it takes a lock and rescans the environment block: the reporter measured 2,431,189 LLAMA_HC_CN_DEBUG calls and 31,204 GGML_CUDA_DISABLE_CONV_FUSION calls per 256 decoded tokens, ~15 ms of extra host time per graph, and a GPU that sat idle waiting on the host (Windows decode ~40 vs ~45 t/s against v16-ebbb18522-r11; Linux unaffected because getenv() is cheap there). Each switch is now resolved once (a static at the call site, or a file-static where the same flag is read from several sites). The same sweep caches the other debug/A-B flags that sit on per-op, per-graph and per-tensor paths: GGML_CUDA_GCDBG, GGML_CUDA_OP_TIMING, GGML_STREAMDBG, GGML_CUDA_MMB_MARK_LOG, the GGML_META_* staging/gather gates, GGML_CUDA_MMQ_J_MAX/_ROUTED, GGML_Q6_COMPACT_J, GGML_CUDA_DISABLE_MMID_512, GGML_PAIR_2X, GGML_CUDA_GDN_CHUNKED(_BF16), GDN_DBG_*, GGML_CUDA_FA_WMMA_256/MAX_HEAD, GGML_CUDA_QSA_SLICES/IDENTITY, the scheduler's GGML_SCHED_* gates (in ggml-backend.cpp), the meta backend's GGML_META_* gates, and GGML_CPU_MOE_OFFLOAD_THREADS. The exact-affecting conv switches fold into block 08's files; everything else rides in block 15 because it is the last block. The values are read at first use, before any graph runs, so behaviour is unchanged. This also fixes the clean-build warnings: the ggml_backend_graph_optimize_params aggregate now sets marks_only/allocs_only, the meta debug prints test tensor->name[0] instead of the always-true array address, the ignored cudaFree()/hipHostRegister() results are cast to void, the unused llama_kv_cache::v_enabled member (shadowed by the constructor parameter) is dropped, and the test probe's unused jmax is removed. Verified on gfx1201 (ROCm 7.14, 3x R9700): same-seed greedy text byte-identical to the pre-amendment build on the 4B (83eec5e9b4f0), clean backend build with zero warnings. Amended 2026-09-29 (PR ggml-org#64, @briansp2020): three bit-exact RDNA4 verify-band wins folded in. 1) FA band: one 64-column block for query widths 5..8. The band ran ncols1 = 4 (32 columns) per block, so a 5..8-token verify launched two column tiles that each streamed the whole KV cache (q8_0 attention ~70 % slower at n_q = 5 than at n_q = 4). A ncols == 64 band row (512 threads, nbatch_combine 32) keeps 4 warps per Q column group, so every column sees the same KV-row split and combine order; it is used for the ncols1 == 4 band from 4096 KV rows. GGML_HIP_FA_BAND_WIDE=0 restores the two-tile launch. gfx1201 q8_0 n_q 5..8: -10..-15 %. 2) ssm_gate_beta_fused_q8_0 gets an ncols template parameter, so the fused Q8_0 alpha/beta projections + softplus/sigmoid chain serves the 2..8-token verify band too (each token keeps the single-token K order, cross-warp reduction and calc_nwarps_weight() selector). -144 launches per 5-token pass. GGML_CUDA_FUSE_GATE_BETA_VERIFY=0 is the off switch. 3) The residual ADD in front of rms_norm_q8_1 (a standalone k_bin_bcast at 2..8 tokens, folded into the mmvq epilogue at one token) is computed inside the norm kernel (rms_norm_q8_1_f32<.., has_add>). -127 launches per 5-token pass. GGML_CUDA_FUSE_ADD_RMS_Q8=0 is the off switch. All three are default-on (maintainer policy) with per-patch kill switches. Independently re-verified here on gfx1201 / ROCm 7.14 / 3x R9700: test-backend-ops -o FLASH_ATTN_EXT 6354/6354; test-logits-width-probe (27B UD-Q4_K_XL, q8_0 KV, RS=from_w) PASS with every W and row0/row1 hash byte-identical with the fusions on vs off and with the wide band on vs off; the 4B same-seed coherence hash is unchanged (83eec5e9b4f0); clean build. The contributor measured -1.4..-2.6 % ms per server verify step; only ROCm 10.0 was measured on their side. Campaign fold (wip/moe-expert-cache, beta2): the CUDA consumer hooks of the decode-side MoE expert cache. Adds the cache-aware fused gate+up+GLU and down folds, the `MUL_MAT_ID` takeover and the cache-band fusion stand-down in ggml-cuda.cu, the backend iface assignment, the `h2d_pin`/`h2d_scratch` staging helpers the gather uses in common.cuh, and the `stage_gather` return-value guard in ggml-backend-meta.cpp. The engine and the kernel slot lookup are in block 13; the generic interface and the scheduler/Meta delegation are in block 06; the gemma4 `-sm tensor` exclusion is in block 14. These pieces ride here because they interleave with block 15's own fusion and staging code. The campaign's bring-up/A-B instrumentation (the CPU routing profiler, the CPU-computes split, the adaptive staging probe, the debug/verify/progress/timing/fail-alloc knobs and the unsafe GGML_META_NOSYNC) was stripped for the promotion candidate. beta5 cleanup (2026-10-02): removed the campaign's leftover env-gated debug/A-B instrumentation (documented as a follow-up at r16). All 25 knobs are gone, with their defaults hard-coded: GGML_META_DEBUG/CNDBG/SSDBG/SS_FIX/REDDBG/VIEWDBG/BUFDBG/SPLITDBG/TOUCHSRC/UPLOADDBG/ NO_2D/SYNCUPLOAD/STAGEDBG/GCDBG/EXECDBG/READIDS/SYNCEACH/GATHERDBG/GATHER_MODE/ PINSRC/PINHOST/PINRING/D2DSPLICE, plus GGML_SCHED_SYNCDBG and GGML_SET_BYTES/SET_NAMES and the TEMP INSTRUMENT blocks they powered (g_ss_*, g_ring_*, g_meta_gc_*, g_gather_*, g_h2d_*, the pool / per-op timers). Behaviour is unchanged on every default path: build warning-free, MUL_MAT_ID 929/929, byte-identity de8be4d0c90c / 15038c19ddc8, width purity, MTP n3 0.75273. Keeps the functional delivery kill-switches (GGML_CUDA_SPLICE_GATHER, GGML_CUDA_GCDBG, GGML_SCHED_STAGE*, etc.). 521 lines removed. Amended 2026-10-04 (issue ggml-org#59 closed; issue ggml-org#60, gfx1100): the qwen4exp prefill policy fix, folded into block 15 and promoted to release v16-84e76d8a2-r34 from the gfx1100 candidate branch wip/issues-59-60 (based on r27). Three parts: - gate qsa3 off on RDNA3_0 by default (GGML_CUDA_QSA3=1 opts in, =0 force-off) and skip the natural F16 K/V packs when the acting backend will not run the packed kernel. The model probes the packed op through ggml_backend_dev_supports_op (qsa3_op_supported), so disabling qsa3 frees its two full F16 cache copies as well; the RDNA3_5/RDNA4 defaults are unchanged. (ggml-org#59, reporter confirmed.) - keep the generic 4-head lightning-indexer prefill score supported on RDNA3_0 and make the fused-vs-chain decision geometry-aware in build_qsa_top_k's use_wmma. Auto selects the fused op on RDNA3_0 once the score exceeds LLAMA_QSA_SCORE_WMMA_MB (default 64 MiB; 0 = always fused) and keeps the faster chain below it; on the reporter's r=4 geometry 64 MiB is n_kv about 64K, which caps the chain's reserve before the 196K prefill's ceiling. The unfused chain's mul_mat+relu and its chunked form now use ggml_relu_inplace, so the chain never holds the score and its relu at once (bit-identical, and it removes the chain's 2x-score reserve peak). LLAMA_QSA_SCORE_WMMA=0/1 forces the chain/fused op. (ggml-org#60.) - account for the FA prefill staging arena in llama_get_memory_breakdown / --fit, bounded by min(2*GGML_CUDA_FA_STAGE_MAX_MB, F16 attention-KV size), so --fit no longer under-provisions a deep prefill (issue ggml-org#33 follow-up). Verified on the r27-based candidate on gfx1100 (ROCm 7.14): clean build; FLASH_ATTN_QSA 23/23 (26/26 with GGML_CUDA_QSA3=1), LIGHTNING_INDEXER 225/225, TOPK_QSA 4/4, FLASH_ATTN_EXT 6354/6354; dense 27B same-seed 1acb04bd9104, byte-identical to r20. Re-folded onto r33 with no conflicts; full record in wip/issues-59-60/ and WORKLOG.md 2026-10-04 (r34). Amended 2026-10-05 (PR ggml-org#91, @briansp2020): BF16 hyper-connection hc_mix fusion (Flash-Next GSQ-RCO). The ISTA-DASLab GSQ-RCO quants keep hc_{attn,ffn}_{down,up,inject} in BF16, and the fused HC_MIX op was Q8_0-only, so those models ran the six-dispatch unfused chain (96 mixers per token, a dispatch-bound decode). The op gains a BF16 arm that replays the unfused BF16 chain in three dispatches with the same per-thread K order and reductions (bit-identical): hc_mix_rms_gamma_quant without the quantize; the down rows and the inject rows in one grid (mul_mat_vec_f<bf16,float,nt,256>); and the up rows of the four hc streams of one column plus the gated collapse in one block (mul_mat_vec_f_vb<bf16,160,4> and dsv4_hc_pre_f32). The CUDA support predicate requires hc_lr == 320 (the emulated mmvf block) and a GPU-resident buffer; LLAMA_HC_MIX_BF16=0 keeps the chain. The CPU HC_MIX reference gained the matching BF16 arm and, with the new BF16 test-backend-ops cases, exposed and fixed a pre-existing dst-stride bug for the no-inject case at nt > 1 (the reference hard-coded the n_embd + hc row stride; it is n_embd without a tail). Gates: test-backend-ops HC_MIX 30/30 (20 Q8_0 + 10 BF16) on gfx1201; fused == unfused byte-identical for nt 1/3/5/8 on a BF16-hc Flash-Next fixture (+3.8 % tg256 on the 4-layer fixture). Record: WORKLOG.md 2026-10-05 (r37), wip/rdna4-hc-mix-bf16/. Amended 2026-10-06 (OPEN 1 / TODO ggml-org#42): the arena slot-count fix, the layer-uniform re-size, the MTP draft cap default, and the cli-only wide-prefill drop. - alloc_table_locked kept t.slots at the REQUESTED count after the r19 slot-count retry shrank the arena, so a shrunk table advertised more slots than its arena held and the decode band read past the allocation (silent corruption: `////` + MTP acceptance 0.007 on an over-filled arena). t.slots is now the achieved count, and g_arena_bytes accounts the achieved count too. - The arena is allocated per LAYER as a unit (alloc_all_locked): all of a layer's gate/up/down tables get the SAME slot count, or an expert resident for one role is not resident for the others and the shared remap names the wrong expert. A shortfall now re-sizes the whole layer down (largest expert slice allocated first, retry at the achieved minimum) instead of shrinking one table; a final non-uniform check disables the cache rather than risk a role mismatch. - MTP draft compute-buffer cap default 512 (MTP_DRAFT_N_UBATCH=0 restores the target's -ub): the draft's encoder-injection buffer is sized for its n_ubatch chunk, so at the target's full -ub it held ~1.6 GiB/device that the MoE arena could use. - The wide-prefill compute-layout drop (TODO ggml-org#42) is now default ON for llama-cli only (common_params::drop_compute_buffers -> llama_context_params -> llama_cparams): llama-cli is a single-wide-prefill workload, where the drop gives a large arena and meets the -ub 8192 decode DoD (measured 78.7 t/s decode / 1683 t/s prefill at 16k, coherent, MTP acc 0.9245). llama-server leaves it off: a later wide prefill needs a contiguous ~12.4-12.9 GB layout back and nothing the arena can yield is large enough (-ub 4096 reclaims, -ub 8192 does not), so a server must not drop. LLAMA_DROP_COMPUTE_BUFFERS overrides either default for A/B. The drop's extra arena reserve now defaults to 0 (the compute-buffer margin covers the small post-drop growth; LLAMA_DROP_EXTRA_RESERVE_MIB=N raises it). A non-uniform arena layer no longer corrupts: the over-filled case degrades one layer to a uniform 3 slots and stays coherent (72.2 t/s).

@briansp2020

Campaign fold (wip/moe-expert-cache, beta2): adds the generic backend expert-cache interface -- `moe_cache_update`, `moe_cache_take_over`, `moe_cache_promote` and `moe_cache_gather` in the backend iface (ggml-backend-impl.h) -- so a backend can drive a residency policy for a host-resident routed-expert table, plus the scheduler half of the consumer: the routed-expert rebalance onto the layer's owning device and the merged per-layer MoE split (ggml-backend.cpp), the device-side expert gather that defers a gatherable routed table from whole-tensor H2D staging, the input-takeover fast path with its deferred promotion pass, the Meta delegation to the per-device simple backends (ggml-backend-meta.cpp), and the `#define GGML_ENV_STR` cache macro the campaign's gates use. The cache engine and its CUDA consumers land in block 13; the gemma4 `-sm tensor` exclusion is in block 14. The CPU and RPC backend vtables name the four new iface fields (NULL) so the set stays warning-free. Amended for the promotion candidate (beta4): keep the scheduler changes prefill-neutral. The device gather now serves only the below-`sched_stage_min_tokens` band (above it the r16/r26 whole-shard staging wins: Q4_K_M ub8192 staging 5283 vs gather 4035; `-sm layer` ncmoe 40 4511 vs 2625), and the routed-expert rebalance is gated to the decode/verify band (`MUL_MAT_ID` `ne[2] <= 8`), where the cache runs. Prefill then matches `main` (r28) at every measured cell while the below-gate gather wins (ub512 +17 %, ub1024 +19 %) are kept. beta5 (2026-10-02): the device gather now also serves a large (>= 224 MiB) host-resident expert table above the width gate. Re-measured with --lazy-mode off (the qwen4exp PLE mmap warm-up confounds every prefill comparison): qwen4exp 450 MiB ub8192 gather 3052 vs staging 1403; Q8_0 272 MiB 3595 vs 3552; Q4_K_M 144 MiB gather 4291 vs staging 5624, so Q4_K_M keeps the width gate. r35 (2026-10-04): issue ggml-org#87's synchronous graph-input fix is now default-on. The r33 `GGML_SCHED_SYNC_GRAPH_INPUTS=1` A/B candidate still reads the variable, but unset means enabled; `=0` restores the r26 async behaviour. The reporter confirmed it clears their recurrent-state H2D race, and same-seed output is unchanged both ways. r5 (2026-10-05, PR ggml-org#96): re-stage a host-resident expert table for a second MUL_MAT_ID consumer in a later split. `ggml_backend_sched_split_graph` registered a weight as a split input only when it first created the copy, so when the same expert weights fed `MUL_MAT_ID` in two splits the later split reused the earlier split's copy. With the expert cache armed that copy was taken over by the 1-row decode-band consumer (`moe_cache_take_over` aliases it and the op never fills it), so the later wide consumer read stale bytes. qwen4exp's unmasked MTP export (a full-row last-layer FFN next to the gathered logits tail) is exactly that second consumer: the wide op produced NaN in `t_h_nextn`, which reached the drafter's KV and collapsed MTP draft acceptance across every slot. When the copy already exists and the node is a `MUL_MAT_ID` reading WEIGHTS through src[0], register the weights as an input of the current split as well (once per split), so they are staged for this split's routing. Single-consumer graphs are unchanged. Contributor PR ggml-org#96 (@briansp2020). Folded 2026-10-07 (TODO ggml-org#45): the stale `--load-mode none` + `-sm tensor` host-expert warning is removed. Its stated mechanism (host experts landing in a pageable CPU_REPACK buffer) no longer holds: the loader pins the routed experts in `ROCm_Host` and the fault no longer reproduces (14/14 clean on r24). Record: archive/work/host-pinned-buffer-crash/.

@briansp2020

The RDNA memory campaign, squashed into one block: six validated wins that remove ~3.4 GiB/GPU of compute-buffer VRAM and ~1.2 GiB of host buffer on qwen4exp (ctx 204800 / ub 2048 / q8_0 KV), plus ~800 MiB/GPU + 800 MiB host on every dense model, at a measured ~1.3 % prefill / ~0.3 % decode cost. Each win keeps an environment kill-switch for A/B (the V4 arm is opt-in). W1 - QSA score-chain memory (GGML_QSA_SCORE_MEM, default 1) relu before the 4-D reshape (elementwise, bit-identical, but it stops the allocator from holding the mul_mat result and its relu simultaneously) + n_blocks-chunked assembly with ggml_concat past a size threshold. qwen4exp ub 2048: 6690.40 -> 4450.40 MiB/GPU. W2 - derived QSA per-block bias + derived visibility (GGML_QSA_DERIVED_BIAS, GGML_QSA_DERIVED_VIS, LLAMA_QSA_SPARSE_FA, defaults 1/1/sparse): the n_blocks x n_tps F32 bias tensor and the additive kq mask are no longer materialised; the fused top-k/FA kernels derive both in-kernel from compact per-cell state. Includes the input-fill null guards (a tensor the graph never consumes is left unallocated by the gallocr, so every fill site must tolerate buffer == nullptr) and the llm_graph_input_attn_k null-mask guard. qwen4exp ub 2048: 4450.40 -> 3251.39 MiB/GPU, host 1262.70 -> 63.69 MiB. W3 - keys-only QSA indexer cache (LLAMA_QSA_KEYS_ONLY, default 1): the indexer scores blocks of keys and never reads a stored value, so its cache is created with no V tensor. Indexer KV 956.26 -> 318.76 MiB/GPU at ctx 204800 (the store is replicated on every GPU). W4 - ggml-alloc: release view sources whose views are never consumed (the counting pass increments view_src->n_views for every view node, but the free pass only decrements it when the view is released - a view with no consumers is never released, so the inflated count blocks the view source's release and in-place reuse). Repro 56.00 -> 16.00 MiB. No gate: it is a bug fix, A/B with beta/block-15-campaign-wins/ab/w4-revert.patch. V3 - derived kq mask for the plain attention path (LLAMA_KQ_MASK_DERIVED, default 1): the packed n_kv x n_tps F16 mask and its host mirror are no longer materialised - the MMA FA kernel derives each cell's visibility from compact per-cell state (GGML_OP_FLASH_ATTN_EXT gains src[5..7]: cell positions, per-token hi/lo bounds). The packed tensor is still created in every graph, so it simply ends up with no consumer: the allocator leaves it unallocated and the existing fill guards skip the fill (no model list, no consumer can be mis-served). The predicate is proven cell-by-cell value-identical (including both SWA cache variants and non-causal) and gates itself off on every unsupported case (alibi, multi-sequence, small/verify batches, non-MMA backends, M-RoPE ext clause). Dense models: -799.20 MiB/GPU and -799.21 MiB host; gemma-4-E4B (ISWA) -809.18/-809.18; gemma-4-31B (ISWA) -811.17/-811.18; scaling exactly n_kv x n_tps x 2 B. The backend support probe is skipped entirely on a multi-stream KV cache (n_seq_max > 1 without kv_unified, and deepseek4, which keeps per-sequence streams even in unified mode): it forces a single-sequence probe graph, and on such a cache deepseek4's indexer stream layout asserts, while the derived form is unreachable there anyway (kq_mask_derivable() requires a single stream). V4 - native q8_0 K/V in the flash-attention kernels (GGML_CUDA_FA_KV_NATIVE, OPT-IN, default 0): the whole-cache F16 staging scratch and its per-ubatch conversion are replaced by dequantising each 16-byte staged chunk while the shared K/V tiles are loaded (a chunk is 8 elements = a quarter q8_0 block, and every MMA/TILE batch is a multiple of 8 elements, so a chunk never straddles a block). The staged values are bit-identical to dequantize_block_q8_0_f16. Dense models: -744 MiB/GPU (4B), -632 (27B), -1224 (gemma-4-31B) at ctx 204800/ub 2048, more at smaller ub; qwen4exp unchanged. Ships opt-in: it costs ~1.7-1.9 % prefill (the lost cp_async pipeline - a quantized source cannot be copied asynchronously) with a large memory win and no cheap way to close the gap; decode is untouched. TILE and VEC already read bf16/f16 natively. Validation (2026-09-10, 3x R9700/RDNA4, ctx 204800, q8_0 KV unless noted; full record: beta/block-15-campaign-wins/): - reserve matrix, all five models x ub 2048/1024/512 x V4 off/on: every number matches the per-win records (the wins compose additively: -2240 W1, -1199 W2, 0/-637.5 W3, -799 V3, -744/-632 V4 per GPU). - same-seed generated text byte-identical on 4B, 27B, gemma-4-E4B (ISWA), gemma-4-31B (ISWA) and qwen4exp across every gate combination (V3 on/off x V4 on/off; qwen4exp W1/W2/W3/W4/V3/V4 flips - 7 configurations), at a short and a 40k-token prompt. - adaptive-MTP gate unchanged: 27B inline draft 0.76744 (66/86, mean 3.28) identical in all four gate combinations; qwen4exp draft 0.44262 (54/122) identical in all six, and equal to the block-14 baseline; MTP still +26 % over plain decode. - test-backend-ops: FLASH_ATTN_EXT on ROCm0 (both V4 gates) and CPU, the derived FA cases, VIEW/CONT/CPY/DUP/CONCAT; test-alloc and test-batch-alloc; the W4 repro goes 56.00 -> 16.00 MiB and the revert patch restores ggml-alloc.c byte-identically. - prefill cost, interleaved same-binary A/B: V3 -1.28 % pp20480/ub 2048 on the 4B (the 27B is neutral, +0.28 %), V4 a further -1.85 % (4B) / -1.72 % (27B); decode within noise in every case. Known pre-existing issue found during this validation (NOT introduced here, reproduces on block 14): gemma-4-E4B-it on 3 GPUs with -sm tensor aborts in the meta splitter (ggml-backend-meta.cpp:1177, a FLASH_ATTN_EXT node whose K source has zero extent on one buffer because n_head_kv = 2 < 3 devices). It works on 1 GPU, on 2 GPUs, and on 3 GPUs with -sm layer. All other models are unaffected. Out of scope for this block. Beta staging, gate table and the full validation record: beta/block-15-campaign-wins/README.md; the per-win designs: wip/arch-independent-memory/{V3-DERIVED-KQ-MASK-PLAN,V4-NATIVE-Q8-KV-PLAN, DERIVED-MASK-DESIGN}.md and wip/qwen4exp/qsa-memory/. Amended 2026-09-10 with V5 - native bf16 K/V in the MMA flash-attention kernel (gated by the SAME GGML_CUDA_FA_KV_NATIVE switch as V4, default 0, i.e. opt-in): a bf16 K/V cache no longer needs the F16 staging scratch. A bf16 source row and the F16 tile it feeds have the same byte layout (2 bytes per element, 16-byte chunks), so the MMA loader converts each staged chunk in registers (bf16 -> f32 -> f16, bit-identical to the launcher's ggml_get_to_fp16_cuda(GGML_TYPE_BF16)) instead of copying it, and the scratch sizing plus the whole-cache conversion pass are skipped for that operand. The staging source is now a single type code (fattn_kv_native_t{FATTN_KV_NATIVE_NONE,Q8_0,BF16}) shared by the launcher, the alloc-size query and both kernels, so they cannot disagree on whether the scratch exists (the old use_q8_K/use_q8_V flags are subsumed). - reserve (ctx 204800, ub 2048, bf16 KV, arm on): 4B 968.86 -> 256.86, 27B 1072.86 -> 488.86, gemma-4-E4B (ISWA) 1062.89 -> 404.89, gemma-4-31B (ISWA) 2068.89 -> 716.89 MiB/GPU - exactly the F16 cache's footprint; ub 1024/512 win 756/778 (4B) and 746 (27B); qwen4exp is unaffected (its FA path never staged bf16 - f16 == bf16 == on/off there), and the TILE/VEC paths already read bf16 natively (ub 8: 8.09 both ways). - same-seed generated text byte-identical (bf16 native vs staged vs f16) on 4B, 27B, gemma-4-E4B and gemma-4-31B, short and 3k/40k-token prompts, with the V3 derived mask active; the two wins are additive (4B: derived mask -799.2 MiB, bf16 scratch -712.0 MiB, both verified independently). - adaptive-MTP gate unchanged: 27B 0.82716 and qwen4exp 0.44262 identical with the arm on and off; the recorded q8_0 references (27B 0.76744, qwen4exp 0.44262) are unchanged by the refactor. - test-backend-ops FLASH_ATTN_EXT: 7859/7859 on ROCm0 and CPU; the 2704 bf16 K/V cases (every head size, including the 576/512 MLA V-is-K-view layout) and the 365 q8_0 cases pass with the arm on and off, with identical case lists. - prefill cost (interleaved same-binary A/B, arm off -> on): the conversion itself is free (native bf16 staging measures within 0.2 % of an F16 cache), but dropping the scratch costs 4B -0.2 % (pp2048), +0.3 % (8192), -1.06 % (20480), -2.36 % (40960) and 27B -0.76 % (20480); decode is within 0.1 % everywhere. The launcher's F16 copy is a dense, normalized copy of the cache view (the GQA heads are interleaved - for the 4B nb[1] is 2048 B for a 512 B row), and the native path re-reads that view on every staging pass. Hence opt-in, exactly like V4 (D9: a sub-2 % loss with a large memory win and no cheap fix ships behind a switch). Mixed K/V cache types (bf16+q8_0, f16+q8_0, either direction) remain a pre-existing trap, documented but not fixed here: the attention falls off the GPU kernel path (graph splits 18 vs 2, ~1.5 GiB host compute buffer, 4B pp2048 7924 -> 640-1049 t/s). Same-type K/V is the practical choice. 2026-09-10 RDNA3_5 (gfx1151) amendment: V3 (derived kq mask) now engages on a ROCm/HIP iGPU. The probe rejected GGML_BACKEND_DEVICE_TYPE_IGPU as "not the CUDA/HIP backend", so on Strix Halo V3 was silently disabled and its ~800 MiB compute + ~800 MiB host win was lost. The device-type gate now accepts IGPU (the ROCm/HIP reg name is still required), and the derived form additionally requires a single KV stream: build_attn_mha derives the attention stream count from the KV tensor (k->ne[3] == the cache's n_stream), so with n_seq_max > 1 -- even for a single-sequence ubatch -- the derived tok_lo was built for one stream while Q was reshaped for n_seq_max streams and the shape assert in ggml_flash_attn_ext_add_kq_derived aborted context creation. kq_mask_derivable() now rejects n_stream != 1, so a multi-slot context keeps the packed mask (no crash) and a single-stream context gets the win. Validated on Strix Halo (gfx1151, ROCm 7.14): V3 on/off byte-identical over 7 KV/arm configs x 16 runs x 129 tokens (2064 gate cells each); reserves reproduce the RDNA4 numbers exactly (4B ctx 204800/ub 2048: V3 -799.20 MiB compute / -799.21 host, V5 bf16 968.86 -> 256.86, V4 q8_0 1001.13 -> 257.13); the derived test-backend-ops cases pass on ROCm0; and llama-server --parallel 4 (the config that aborted before) now serves and passes the 16-run gate. Discrete GPUs are unaffected (they report GGML_BACKEND_DEVICE_TYPE_GPU and normally have n_stream == 1). Amended 2026-09-11 (9th re-cut): fix a variable-shadowing bug in build_attn_qsa that broke the LLAMA_QSA_SPARSE_FA=0 dense arm for every model/type. The V2/V3 refactor wrapped the top-k mask chain in `if (kq_mask != nullptr) { ... }` and added an outer `kq_mask_top_k` declaration, leaving the chain's own `ggml_tensor * kq_mask_top_k = ggml_set_rows(...)` inside the block as a NEW local. The chain was therefore built (whenever the mask existed) but its result never reached the attention: build_attn_mha got the outer `nullptr`, the chain's nodes were unreachable from the graph output (so ggml never emitted them), the packed mask lost its only consumer (so the allocator dropped it and set_input_kq_mask was skipped), and the dense arm attended with no mask at all - a full causal leak. Observed as qwen4exp perplexity ~1.02-1.06 on every KV type (delivery: 6.49-6.55), and as ~1.03 on RANDOM text (delivery ~18.4), with dense-arm greedy text differing from the delivery. Fixed by assigning the block's result to the outer variable (dropping the inner `ggml_tensor *`). The production (sparse) arm never enters that block, so it is bit-identical before and after. See beta/block-15-campaign-wins/BETA-TESTING.md 4c and WORKLOG.md 2026-09-11 (11). 2026-09-14 amendment (issue ggml-org#30, "wider configuration"): V4 is no longer opt-in for the sub-F16 quantized cache types, and q4_0 gains a native arm. - GGML_CUDA_FA_KV_NATIVE is now a three-state policy: unset = auto (native q8_0/q4_0 ON, native bf16 OFF), =1 forces both on, =0 forces the pre-amendment F16-staging path (the escape hatch). The whole-cache F16 staging pass is proportional to n_kv and runs on every decode step, so removing it for the sub-F16 quants is worth +23 % tg64 at d65536 on gfx1201 (q8_0 18.92 -> 23.29 vs stock 22.43) at ~1.2 % prefill, and it recovers ~744 MiB/GPU of scratch. - Native q4_0 arm beside q8_0/bf16 (ggml_cuda_fattn_dequantize_q4_0_chunk, arithmetic-identical to convert.cu's dequantize_block_q4_0) in the tile and MMA kernels: q4_0 d65536 19.72 -> 22.82 vs stock 21.03, ~1.3 % prefill. - Bit-identical to the staging path (same-seed greedy text) and width-pure (W = 1..8 one logits hash per KV type); MTP n_max 3 acceptance unchanged (0.75182). Full evidence: wip/issue-30-mtp-decode-regression/MEASUREMENTS.md. - Side effect: also fixes the adaptive-MTP high-context load failure (--spec-draft-n-max 12 -c 196608 q8_0, default 4 slots): the F16 staging scratch was exactly the 260 MiB the draft context was short, so it now loads and generates. 2026-09-15 amendment (issue ggml-org#30, the reporter's r3 review): the q4_0 NaN fix, the prefill band split, and native arms for the last four quantized K/V types. - Mixed-K/V kernel contract (the reporter's 4 NaN failures in test-backend-ops -o FLASH_ATTN_EXT). The tile kernel is instantiated with ONE type_KV for both operands and builds its native descriptors from that compile-time type, while launch_fattn chose its native read per tensor; a mixed pair (K=q4_0, V=f16) therefore fell back to the F16 tile with K's staging skipped, and the kernel read raw q4_0 as F16. launch_fattn now takes the kernel's native type explicitly (kv_native_kernel; the tile passes its type_KV, the vec passes NONE, the MMA keeps per-operand semantics). llama.cpp rejects mixed caches, so only the op test could see it. - ggml_cuda_flash_attn_ext_get_alloc_size's TILE case never learned the q4_0 arm, so a q4_0 cache kept reserving the F16 scratch the launcher no longer used: -c 196608 q4_0 849.04 -> 123.04 MiB -- the arm's memory win had never actually been delivered. - Prefill band split: a prefill (n_q > 8) stages the operand into F16 while decode/verify (n_q <= 8) keeps the native read, recovering the quantized-KV prefill cost (q8_0 pp150000 661.0 -> 691.4 on 1 GPU, 996.0 -> 1076.9 on 2, 1111.4 -> 1199.0 on 3; q4_0 694.5). The staging scratch comes from a new per-context, per-stream arena (ggml_backend_cuda_context::fattn_stage / fattn_stage_get) instead of the compute-graph reserve, which sized it for n_ctx (~726 MiB at a 200k context) even though the real scratch tracks the prefix; a multi-token graph is never CUDA-graph captured, so a runtime allocation there is safe. Bounded by GGML_CUDA_FA_STAGE_MAX_MB (default 512 MiB, 0 = unbounded). Arch-gated: RDNA3_5 keeps the native read at prefill (it wins at every depth there: gfx1151 9B q8_0 pp16k 1410 vs 1405, pp65k 1059 vs 1043), RDNA4/RDNA3_0 stage. - Native arms for q4_1 / q5_0 / q5_1 / iq4_nl (ggml_cuda_fattn_dequantize_{q4_1,q5_0,q5_1, iq4_nl}_chunk, arithmetic-identical to convert.cu), closing the last gap in the V4 set -- the four types stage through F16 otherwise. tg64 @ d32768 staged -> native: gfx1201 +9.9/+10.8/+12.5/+8.8 % (q4_1 23.14 -> 25.44, q5_0 22.16 -> 24.56, q5_1 22.23 -> 25.00, iq4_nl 22.92 -> 24.94) with prefill unchanged; gfx1151 (9B) +22.9/+26.0/+26.7/+22.0 % for a 0.6-1.1 % prefill cost. Two bugs the op test caught on the way: the tile loader chose the native branch with a hand-written `type_KV == Q8_0 || Q4_0` test (the new instantiations then took the F16 branch and read an unwritten staging buffer -- the hsk=72 NaNs), and the q5_0/q5_1 chunk helpers took the low nibble in both halves. - Same-seed greedy text is bit-identical native vs staging for all eight KV types on both arches, and test-backend-ops -o FLASH_ATTN_EXT is 5951/5951 on both. Width purity: gfx1151 pure for all eight types; gfx1201 pure except the documented pre-existing q4_0 logits-level band edges (GREEDY-PURITY.md 36, whose guarantee is now stated at the text/acceptance level). 2026-09-15 amendment (build time): the tile kernel's native-KV type axis is instantiated in the generated template-instance TUs again instead of implicitly in the tile dispatch TU. Block 03 turned type_KV into a template parameter of ggml_cuda_flash_attn_ext_tile_case so the tile kernel could read BF16 (and later the quantized types) natively, while DECL_FATTN_TILE_CASE / EXTERN_DECL_FATTN_TILE_CASES kept covering only F16/BF16. Because the dispatch in fattn-tile.cu has an unconditional `case` per native type, the other six types were then instantiated *implicitly in that TU*: 12 head-size combinations x 6 types = 72 case instantiations (96 symbols including the 24 externed F16/BF16), each pulling in both softcap variants and the whole ncols2 chain. With block 08's FA_QUANTS default (8 diagonal pairs) that single TU took 509 s of a 538 s `-j16` backend build on a 16-core machine (measured 2026-09-15, gfx1201) -- it *was* the critical path, and no amount of build parallelism could help. The macros now cover every type the dispatch can select, so the 12 generated template-instances/fattn-tile-instance-*.cu files instantiate 8 cases each and the dispatch TU holds only `extern` declarations. Build time: `ggml-hip` 538 s -> 330 s, and the tile dispatch TU itself drops from 509 s to under 10 s. The remaining critical path is the `fattn-mma-f16` instance set, and that one is ours too: this block's native-KV arms instantiate the whole WMMA kernel once per KV type *inside every instance TU* (the arm chain in `ggml_cuda_flash_attn_ext_mma_f16_case`), so each of the 21 instance files compiles 8 kernel copies. Measured on the same TU (`ncols1_4-ncols2_4`): object 0.90 MB -> 7.26 MB, compile 6.7 s -> 229 s. Fixing it needs either a finer generated-file granularity (~21 -> ~150 TUs) or a runtime KV-type dispatch in the WMMA loader, i.e. a code-path change with its own A/B, so it is left as a follow-up rather than folded in here. Nothing about the kernels changed -- same template arguments, same flags, same device code; only the TU that emits them differs. Verified: `test-backend-ops -o FLASH_ATTN_EXT` 5951/5951 with zero failures, the 27B prose text hashes bit-identical to the pre-amendment build (q8_0 472b282950b5, q4_0 118eb7f5fe85, f16 70960317a203), and `tg64@32768`/`pp8192` reproduce the recorded values within run-to-run noise across five KV types. Amended 2026-09-17: the HIP build now passes `-Wno-pass-failed`. The FA kernels carry bare `#pragma unroll` hints on runtime-bounded loops, so AMDGPU clang reports every one it cannot honour; the native-KV arms instantiate the WMMA kernel per KV type per instance TU and multiply that into ~10k "loop not unrolled" warning lines per clean build. The hints are advisory and the pass already failed, so the build output is unchanged -- only the flood is gone. HIP-only (the diagnostic is Clang/AMDGPU). See wip/build-time-regression/. Amended 2026-09-18: the tile template instances are generated per (head size, KV type) instead of per head size, matching the per-type axis this block introduced. Build-time only: the same 96 instantiations, identical symbols. Amended 2026-09-25 (issue ggml-org#45): the RDNA4 GQA-6 decode/verify flash-attention band. At head 256 with GQA 6, the tile kernel can only fold ncols2 = 2 (its ncols2 must divide the GQA ratio), so the whole n_q <= 8 band fetched and dequantized every K/V element once per head pair -- three times per query row at 276 GB/s on q8_0. The band now routes the whole band (n_q = 1 included) to the WMMA kernel with the GQA group folded into ncols2 = 8 (existing (256, ncols1 4|2, ncols2 8) instances, no new instantiations): one block covers the group, and the KV is split round-robin over a fixed P = nsm blocks per output tile (launch_fattn, GGML_HIP_FA_BAND_WMMA_SPLIT override). P depends on neither n_q nor the KV length, so decode and every verify width accumulate the same partials in the same order -- the round-robin form is what keeps the GREEDY-PURITY band invariant (a contiguous stream-k split moved its split points with the 256-padded KV length and broke it). - kernel chooser, ncols dispatcher and launch_fattn share one self-contained predicate (ggml_cuda_fattn_band_wmma_applies), and launch_fattn also requires the ncols2 == 8 template parameter, so a disagreement can never launch the 2-D grid into a kernel whose fast path is compiled out. - coverage follows ggml_cuda_fattn_kv_native_type, not a hand-written type list: every native quantized K/V type wins (q8_0, q4_0, q4_1, q5_0, q5_1, iq4_nl); f16 stays on tile (DRAM-bound) and bf16 stays on tile while its native arm is opt-in. - default ON (maintainer policy); GGML_HIP_FA_BAND_WMMA=0 opts out (=2 picks ncols1 = 2). GGML_CUDA_FA_WMMA_256=0 / GGML_CUDA_FA_WMMA_MAX_HEAD<256 still disable the band. Measured on gfx1201 (R9700), 27B qwen35 head 256 GQA 6, q8_0 KV, op level kv 16384 (tile -> band, us/run): n_q 3 326 -> 150, n_q 5 524 -> 225, n_q 8 814 -> 230; n_q 1 141 -> 148 (q4_0/q4_1/iq4_nl are faster at n_q 1 too). End to end, 27B UD-Q4_K_XL draft-mtp n3 at ~40k ctx, band off -> on: 1 GPU 43.3 -> 51.3, 2 GPU tensor 67.9 -> 76.6, 3 GPU tensor 77.6 -> 88.8 t/s (+13-18 %), while pp512/2048/4096 are flat (the band is gated to n_q <= 8, so prefill is untouched by construction). Purity re-verified with the band on: test-backend-ops -o FLASH_ATTN_EXT 6340/6340 on ROCm0 (389 GQA-6 cases across all eight KV types, both layouts); plain == draft-mtp byte-identical on all eight KV types at ~40k ctx, 3-GPU tensor (acceptance 0.83-0.86), and q8_0/q4_0 text identical across --spec-draft-n-max 3/5/7 (verify widths 4/6/8); 1 GPU layer, 2/3 GPU tensor and 3 GPU layer are all pure. New qwen35-band eval and perf cases added to tests/test-backend-ops.cpp. Folds beta/mmb-general: 0002 (qsa3 packed WMMA), 0003 (qwen4exp always-QSA flip), 0004 (HC16 native-BF16 producers), 0005 (fused indexer top-k), 0008 (HC16/beneficial defaults on ggml-cuda.cu + llama-context.cpp), 0011 (qsa3 on RDNA3_0), 0013 (hc_gate_mix), 0016 (-lzm auto + managed PLE gate), 0017 (HC BF16 streams), 0018 (out_xn BF16), 0019 (indexer relu+head-sum), 0020 (F32 activations on eval callback), 0021 (sparse MTP-draft attention), 0022 (derived indexer cache ON), 0023 (QSA gfx1151 crossover), 0024 (HC16 per-backend state), 0026 (meta child graph_optimize markings). Amended 2026-09-27 (issue ggml-org#53): the derived kq-mask inputs must not be read on the host, and the frontend test cases must set tensors through the backend API. The r10 fully-masked KV-group skip built its batch-wide bitmap by dereferencing `tok_lo`/`tok_hi` on the CPU in `launch_fattn`. Those are host graph inputs (GGML_TENSOR_FLAG_INPUT), and the backend scheduler copies such an input to the compute backend, so the tensors the launcher sees are the *device* copies. A host read of device memory is an access violation wherever the allocation is not CPU-mapped - on Windows/WDDM it is a 0xC0000005 in ggml-hip.dll on the very first prefill ubatch (the reporter's r10 crash) - and a race against the in-flight copy everywhere else. The batch-wide [lo_min, hi_max] window is now computed inside `flash_attn_kq_derived_blocks` from the device-resident `tok_lo`/`tok_hi`; the bitmap it produces is unchanged (same integer min/max), so the skip is still exact. Verified on gfx1201: test-backend-ops -o FLASH_ATTN_EXT 6354/6354 with the MMA and the tile kernel (and the 20 derived/mask_hole cases specifically), FLASH_ATTN_QSA 26/26, and 4B Q8_0 same-seed text identical across {skip on, skip=0, derived=0}. The same class of bug is fixed in tests/test-backend-ops.cpp: the derived (`cell_pos`/`tok_lo`/`tok_hi`), packed `mask_hole` and `FLASH_ATTN_QSA` (`idx`/`m`) initializers wrote `t->data` directly instead of going through `ggml_backend_tensor_set`, so on Windows the first `derived=1` case segfaulted before any derived case ran (also reported in issue ggml-org#53). The QSA tensors are block-14 test code; the fix rides in block 15 because it is the last block. Amended 2026-09-29 (issue ggml-org#65): cache the getenv() lookups on the fusion and staging hot paths. ggml_can_fuse_subgraph_ext() read LLAMA_HC_CN_DEBUG for every candidate fusion window (millions of calls per pass), ggml_cuda_try_fuse() read it again in the four hc_combine_norm matcher paths, and gdn_conv_enabled()/ple_conv_enabled() read GGML_CUDA_DISABLE_CONV_FUSION once per conv launch. On Linux getenv() is cheap, but on Windows it takes a lock and rescans the environment block: the reporter measured 2,431,189 LLAMA_HC_CN_DEBUG calls and 31,204 GGML_CUDA_DISABLE_CONV_FUSION calls per 256 decoded tokens, ~15 ms of extra host time per graph, and a GPU that sat idle waiting on the host (Windows decode ~40 vs ~45 t/s against v16-ebbb18522-r11; Linux unaffected because getenv() is cheap there). Each switch is now resolved once (a static at the call site, or a file-static where the same flag is read from several sites). The same sweep caches the other debug/A-B flags that sit on per-op, per-graph and per-tensor paths: GGML_CUDA_GCDBG, GGML_CUDA_OP_TIMING, GGML_STREAMDBG, GGML_CUDA_MMB_MARK_LOG, the GGML_META_* staging/gather gates, GGML_CUDA_MMQ_J_MAX/_ROUTED, GGML_Q6_COMPACT_J, GGML_CUDA_DISABLE_MMID_512, GGML_PAIR_2X, GGML_CUDA_GDN_CHUNKED(_BF16), GDN_DBG_*, GGML_CUDA_FA_WMMA_256/MAX_HEAD, GGML_CUDA_QSA_SLICES/IDENTITY, the scheduler's GGML_SCHED_* gates (in ggml-backend.cpp), the meta backend's GGML_META_* gates, and GGML_CPU_MOE_OFFLOAD_THREADS. The exact-affecting conv switches fold into block 08's files; everything else rides in block 15 because it is the last block. The values are read at first use, before any graph runs, so behaviour is unchanged. This also fixes the clean-build warnings: the ggml_backend_graph_optimize_params aggregate now sets marks_only/allocs_only, the meta debug prints test tensor->name[0] instead of the always-true array address, the ignored cudaFree()/hipHostRegister() results are cast to void, the unused llama_kv_cache::v_enabled member (shadowed by the constructor parameter) is dropped, and the test probe's unused jmax is removed. Verified on gfx1201 (ROCm 7.14, 3x R9700): same-seed greedy text byte-identical to the pre-amendment build on the 4B (83eec5e9b4f0), clean backend build with zero warnings. Amended 2026-09-29 (PR ggml-org#64, @briansp2020): three bit-exact RDNA4 verify-band wins folded in. 1) FA band: one 64-column block for query widths 5..8. The band ran ncols1 = 4 (32 columns) per block, so a 5..8-token verify launched two column tiles that each streamed the whole KV cache (q8_0 attention ~70 % slower at n_q = 5 than at n_q = 4). A ncols == 64 band row (512 threads, nbatch_combine 32) keeps 4 warps per Q column group, so every column sees the same KV-row split and combine order; it is used for the ncols1 == 4 band from 4096 KV rows. GGML_HIP_FA_BAND_WIDE=0 restores the two-tile launch. gfx1201 q8_0 n_q 5..8: -10..-15 %. 2) ssm_gate_beta_fused_q8_0 gets an ncols template parameter, so the fused Q8_0 alpha/beta projections + softplus/sigmoid chain serves the 2..8-token verify band too (each token keeps the single-token K order, cross-warp reduction and calc_nwarps_weight() selector). -144 launches per 5-token pass. GGML_CUDA_FUSE_GATE_BETA_VERIFY=0 is the off switch. 3) The residual ADD in front of rms_norm_q8_1 (a standalone k_bin_bcast at 2..8 tokens, folded into the mmvq epilogue at one token) is computed inside the norm kernel (rms_norm_q8_1_f32<.., has_add>). -127 launches per 5-token pass. GGML_CUDA_FUSE_ADD_RMS_Q8=0 is the off switch. All three are default-on (maintainer policy) with per-patch kill switches. Independently re-verified here on gfx1201 / ROCm 7.14 / 3x R9700: test-backend-ops -o FLASH_ATTN_EXT 6354/6354; test-logits-width-probe (27B UD-Q4_K_XL, q8_0 KV, RS=from_w) PASS with every W and row0/row1 hash byte-identical with the fusions on vs off and with the wide band on vs off; the 4B same-seed coherence hash is unchanged (83eec5e9b4f0); clean build. The contributor measured -1.4..-2.6 % ms per server verify step; only ROCm 10.0 was measured on their side. Campaign fold (wip/moe-expert-cache, beta2): the CUDA consumer hooks of the decode-side MoE expert cache. Adds the cache-aware fused gate+up+GLU and down folds, the `MUL_MAT_ID` takeover and the cache-band fusion stand-down in ggml-cuda.cu, the backend iface assignment, the `h2d_pin`/`h2d_scratch` staging helpers the gather uses in common.cuh, and the `stage_gather` return-value guard in ggml-backend-meta.cpp. The engine and the kernel slot lookup are in block 13; the generic interface and the scheduler/Meta delegation are in block 06; the gemma4 `-sm tensor` exclusion is in block 14. These pieces ride here because they interleave with block 15's own fusion and staging code. The campaign's bring-up/A-B instrumentation (the CPU routing profiler, the CPU-computes split, the adaptive staging probe, the debug/verify/progress/timing/fail-alloc knobs and the unsafe GGML_META_NOSYNC) was stripped for the promotion candidate. beta5 cleanup (2026-10-02): removed the campaign's leftover env-gated debug/A-B instrumentation (documented as a follow-up at r16). All 25 knobs are gone, with their defaults hard-coded: GGML_META_DEBUG/CNDBG/SSDBG/SS_FIX/REDDBG/VIEWDBG/BUFDBG/SPLITDBG/TOUCHSRC/UPLOADDBG/ NO_2D/SYNCUPLOAD/STAGEDBG/GCDBG/EXECDBG/READIDS/SYNCEACH/GATHERDBG/GATHER_MODE/ PINSRC/PINHOST/PINRING/D2DSPLICE, plus GGML_SCHED_SYNCDBG and GGML_SET_BYTES/SET_NAMES and the TEMP INSTRUMENT blocks they powered (g_ss_*, g_ring_*, g_meta_gc_*, g_gather_*, g_h2d_*, the pool / per-op timers). Behaviour is unchanged on every default path: build warning-free, MUL_MAT_ID 929/929, byte-identity de8be4d0c90c / 15038c19ddc8, width purity, MTP n3 0.75273. Keeps the functional delivery kill-switches (GGML_CUDA_SPLICE_GATHER, GGML_CUDA_GCDBG, GGML_SCHED_STAGE*, etc.). 521 lines removed. Amended 2026-10-04 (issue ggml-org#59 closed; issue ggml-org#60, gfx1100): the qwen4exp prefill policy fix, folded into block 15 and promoted to release v16-84e76d8a2-r34 from the gfx1100 candidate branch wip/issues-59-60 (based on r27). Three parts: - gate qsa3 off on RDNA3_0 by default (GGML_CUDA_QSA3=1 opts in, =0 force-off) and skip the natural F16 K/V packs when the acting backend will not run the packed kernel. The model probes the packed op through ggml_backend_dev_supports_op (qsa3_op_supported), so disabling qsa3 frees its two full F16 cache copies as well; the RDNA3_5/RDNA4 defaults are unchanged. (ggml-org#59, reporter confirmed.) - keep the generic 4-head lightning-indexer prefill score supported on RDNA3_0 and make the fused-vs-chain decision geometry-aware in build_qsa_top_k's use_wmma. Auto selects the fused op on RDNA3_0 once the score exceeds LLAMA_QSA_SCORE_WMMA_MB (default 64 MiB; 0 = always fused) and keeps the faster chain below it; on the reporter's r=4 geometry 64 MiB is n_kv about 64K, which caps the chain's reserve before the 196K prefill's ceiling. The unfused chain's mul_mat+relu and its chunked form now use ggml_relu_inplace, so the chain never holds the score and its relu at once (bit-identical, and it removes the chain's 2x-score reserve peak). LLAMA_QSA_SCORE_WMMA=0/1 forces the chain/fused op. (ggml-org#60.) - account for the FA prefill staging arena in llama_get_memory_breakdown / --fit, bounded by min(2*GGML_CUDA_FA_STAGE_MAX_MB, F16 attention-KV size), so --fit no longer under-provisions a deep prefill (issue ggml-org#33 follow-up). Verified on the r27-based candidate on gfx1100 (ROCm 7.14): clean build; FLASH_ATTN_QSA 23/23 (26/26 with GGML_CUDA_QSA3=1), LIGHTNING_INDEXER 225/225, TOPK_QSA 4/4, FLASH_ATTN_EXT 6354/6354; dense 27B same-seed 1acb04bd9104, byte-identical to r20. Re-folded onto r33 with no conflicts; full record in wip/issues-59-60/ and WORKLOG.md 2026-10-04 (r34). Amended 2026-10-05 (PR ggml-org#91, @briansp2020): BF16 hyper-connection hc_mix fusion (Flash-Next GSQ-RCO). The ISTA-DASLab GSQ-RCO quants keep hc_{attn,ffn}_{down,up,inject} in BF16, and the fused HC_MIX op was Q8_0-only, so those models ran the six-dispatch unfused chain (96 mixers per token, a dispatch-bound decode). The op gains a BF16 arm that replays the unfused BF16 chain in three dispatches with the same per-thread K order and reductions (bit-identical): hc_mix_rms_gamma_quant without the quantize; the down rows and the inject rows in one grid (mul_mat_vec_f<bf16,float,nt,256>); and the up rows of the four hc streams of one column plus the gated collapse in one block (mul_mat_vec_f_vb<bf16,160,4> and dsv4_hc_pre_f32). The CUDA support predicate requires hc_lr == 320 (the emulated mmvf block) and a GPU-resident buffer; LLAMA_HC_MIX_BF16=0 keeps the chain. The CPU HC_MIX reference gained the matching BF16 arm and, with the new BF16 test-backend-ops cases, exposed and fixed a pre-existing dst-stride bug for the no-inject case at nt > 1 (the reference hard-coded the n_embd + hc row stride; it is n_embd without a tail). Gates: test-backend-ops HC_MIX 30/30 (20 Q8_0 + 10 BF16) on gfx1201; fused == unfused byte-identical for nt 1/3/5/8 on a BF16-hc Flash-Next fixture (+3.8 % tg256 on the 4-layer fixture). Record: WORKLOG.md 2026-10-05 (r37), wip/rdna4-hc-mix-bf16/. Amended 2026-10-06 (OPEN 1 / TODO ggml-org#42): the arena slot-count fix, the layer-uniform re-size, the MTP draft cap default, and the cli-only wide-prefill drop. - alloc_table_locked kept t.slots at the REQUESTED count after the r19 slot-count retry shrank the arena, so a shrunk table advertised more slots than its arena held and the decode band read past the allocation (silent corruption: `////` + MTP acceptance 0.007 on an over-filled arena). t.slots is now the achieved count, and g_arena_bytes accounts the achieved count too. - The arena is allocated per LAYER as a unit (alloc_all_locked): all of a layer's gate/up/down tables get the SAME slot count, or an expert resident for one role is not resident for the others and the shared remap names the wrong expert. A shortfall now re-sizes the whole layer down (largest expert slice allocated first, retry at the achieved minimum) instead of shrinking one table; a final non-uniform check disables the cache rather than risk a role mismatch. - MTP draft compute-buffer cap default 512 (MTP_DRAFT_N_UBATCH=0 restores the target's -ub): the draft's encoder-injection buffer is sized for its n_ubatch chunk, so at the target's full -ub it held ~1.6 GiB/device that the MoE arena could use. - The wide-prefill compute-layout drop (TODO ggml-org#42) is now default ON for llama-cli only (common_params::drop_compute_buffers -> llama_context_params -> llama_cparams): llama-cli is a single-wide-prefill workload, where the drop gives a large arena and meets the -ub 8192 decode DoD (measured 78.7 t/s decode / 1683 t/s prefill at 16k, coherent, MTP acc 0.9245). llama-server leaves it off: a later wide prefill needs a contiguous ~12.4-12.9 GB layout back and nothing the arena can yield is large enough (-ub 4096 reclaims, -ub 8192 does not), so a server must not drop. LLAMA_DROP_COMPUTE_BUFFERS overrides either default for A/B. The drop's extra arena reserve now defaults to 0 (the compute-buffer margin covers the small post-drop growth; LLAMA_DROP_EXTRA_RESERVE_MIB=N raises it). A non-uniform arena layer no longer corrupts: the over-filled case degrades one layer to a uniform 3 slots and stays coherent (72.2 t/s).

Fix typo in README

2fed42d

prusnak approved these changes Mar 12, 2023

View reviewed changes

ggerganov merged commit 9661954 into ggml-org:master Mar 12, 2023

Bearsaerker mentioned this pull request Mar 12, 2025

Eval bug: Gemma 3 extremly slow prompt processing when using quantized kv cache. #12352

Closed

jesusmb1995 pushed a commit to jesusmb1995/llama.cpp that referenced this pull request Oct 30, 2025

Merge pull request ggml-org#45 from jpgaribotti/QVAC-6569

3b10875

QVAC-6569: Redirected performance logging to GGML_LOG

wine99 pushed a commit to wine99/llama.cpp that referenced this pull request Feb 27, 2026

Merge pull request ggml-org#45 from cavusmustafa/tmp_fix_multithread

cb92f77

Temp fix for multithreading bug

Seunghhon pushed a commit to Seunghhon/llama.cpp that referenced this pull request Apr 26, 2026

Fix typo in README (ggml-org#45)

8de246c

zhsy12345689 pushed a commit to zhsy12345689/llama.cpp-omni that referenced this pull request Jul 30, 2026

Merge pull request ggml-org#45 from oreomaker/voxcpm

8683781

Support VoxCPM 0.5B, 1.5, 2 inference on llama.cpp

Navigation Menu

Repository navigation

Fix typo in README - #45

Fix typo in README#45
ggerganov merged 1 commit into
ggml-org:masterfrom
marckohlbrugge:patch-1

marckohlbrugge commented Mar 12, 2023

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Conversation

marckohlbrugge commented Mar 12, 2023

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants