A patch collection that brings AMD RDNA-specific performance work to llama.cpp: MTP decode, chunked gated-delta-net prefill, BF16 KV and WMMA flash-attention, fused MoE and k-quant decode paths, a hybrid all-reduce, qwen4exp (Qwen3.8-Flash-Next) support, and an attention-memory campaign that frees several GiB of VRAM.
It ships as 16 patches (block 00 + blocks 01-15) for a clean llama.cpp
checkout at the fork point 84e76d8a2 (upstream master, 2026-09-24
re-base). Each block is a self-contained git am commit, so you can apply
the whole set or pick the ones you want. The mmb (bf16-WMMA weight GEMM) / QSA / indexer
campaign, formerly the 28-patch opt-in archive/work/mmb-general/ set, is now folded into the delivery
blocks — the mmb core into block 08, the catch-all system-operations fixes into block 06, and the
qwen4exp/QSA/HC/indexer work into block 15 — so the 16 patches alone reproduce the full campaign
tree 24bb0f5acb…. archive/work/mmb-general/ is retained only as the historical verification record;
see The mmb campaign is in the delivery. (This is the
state on the beta-integration branch; main still carries r7 + the separate beta set.)
git clone https://github.com/ggml-org/llama.cpp && cd llama.cpp
git checkout 84e76d8a2
bash <path-to-this-repo>/scripts/apply-all.sh . # creates branch rdna-boosts- One-line summary of each block: The 16 blocks
- The folded
mmb/QSA campaign: Themmbcampaign is in the delivery - Apply details, env knobs, server config:
patches/README.md - What changed recently:
WORKLOG.md - Current status and validation: Current state
The
mmb/QSA/indexer campaign is folded into the delivery (2026-09-25, releaser8). The 16 delivery patches now absorb the 28archive/work/mmb-general/patches; the applied tree is24bb0f5acb…and the build is clean on gfx1201. The working plan, per-patch mapping and validation record are inarchive/work/beta-integration/integration.md.
Frozen deliveries are published as GitHub Releases and tagged in this repo
(the tag is the release identity: v16-<fork-point>-r<N>, e.g.
v16-84e76d8a2-r7, where r1 is the re-base, r2 the block-10 MoE-VDR arch-scope fix, r3 the
block-14 hc_combine CPU-reference fix (issue #44) + the beta re-base, r4 the block-15 RDNA4
GQA-6 decode/verify flash-attention band (issue #45), r5 the block-15 f16/bf16 band coverage
(issue #45 follow-up), r6 the block-15 bf16 native default flip, r7 the block-14 Meta-tensor-split
scheduler race fix, r8 the archive/work/mmb-general fold into the 16 blocks (the campaign is now
part of the delivery, no separate beta apply step), r9 the block-15 typed non-swizzled K/V store
fix for the MMA FA prefill loader (issue #47), r10 the block-15 fully-masked KV-group skip that
stops a --kv-unified concurrent prefill from paying for the other slots' cells (issue #48), r11 the
block-13 MoE MMVQ rpb mis-launch fix behind the MUL_MAT_ID backend-ops failure, r12 the block-06
op-offload H2D staging ring (issue #50, merging PR #51 by @briansp2020) with the -sm tensor
op-offload fix it needed, plus the block-06 rename to general system-operations bucket, r13 the
block-06 tiny-CPU-graph heuristic fix that counts the tensors a node reads, so a CPU-offloaded FFN chunk
is no longer serialized on one thread (issue #52), r14 the block-15 fix that computes the derived
kq-mask window on the device instead of reading the device copy of tok_lo/tok_hi from the host (the
Windows 0xC0000005 on the first prefill ubatch, issue #53), r15 the block-06 fix that keeps
host-resident MoE expert weights (MUL_MAT_ID) pinned instead of downgrading them to the pageable mmap;
each later release on the same base
increments N). release.json.release must equal the tag — CI
checks it — and only a tag push cuts a release. Each release carries
rdna-boosts-all.patch, patches.tar.gz, release.json
and SHA256SUMS, so a consumer can pin a tag and verify the artifacts instead
of tracking a moving main.
release.json is the delivery's single source of truth (fork point, canonical
tip/tree, block count, per-artifact sha256); scripts/apply-all.sh,
scripts/validate-set.sh and CI all read it. The container pipeline is
tag-driven, so ordinary commits to main (docs / benchmarks / WORKLOG.md)
only run the cheap patch validation — see CONTAINERS.md for
the release process and the prebuilt ROCm images.
The set targets the RDNA3 / RDNA3.5 / RDNA4 GPU families:
| family | arches | example parts |
|---|---|---|
| RDNA 3 | gfx1100 |
RX 7900 XTX/XT, RX 7800 XT, ... |
| RDNA 3.5 | gfx1150/gfx1151 |
Strix Point / Strix Halo APUs |
| RDNA 4 | gfx1200/gfx1201 |
RX 9060 XT; RX 9070 / 9070 XT |
RDNA4 (gfx120x) sees the most benefit — the WMMA flash-attn path, the chunked-GDN kernel, the k-quant VDR boosts and block 12's internal all-reduce were all first built and validated there. As much of that work as possible is back-ported to the RDNA3/3.5 families instead of being gated off:
- block 02's chunked gated-delta-net bf16/WMMA prefill ships as two
arch-segregated kernels: a dedicated first-gen WMMA port for gfx11
(
gated_delta_net_chunked_bf16_gfx11.cu) next to the RDNA4 kernel; - block 04's WMMA flash-attn is not RDNA4-only despite the block name — RDNA3.0 runs it with the same 576-head limit as RDNA4, RDNA3.5 with a tuned 320-head limit;
- block 10 adds a dedicated RDNA3.5 mmvq parameter table (previously folded into the RDNA2 fallback) on top of the RDNA4 k-quant boosts.
Arch selection is runtime everywhere in the set (device cc /
gcnArchName; there is no compile-time arch gating), so a multi-arch
build such as GPU_TARGETS="gfx1100;gfx1151;gfx1201" yields one binary
that picks the right path on whichever of these it runs on. The one
genuine exception is block 12 — its internal all-reduce is RDNA4-only
(gfx1200/gfx1201) and falls back to RCCL elsewhere (see
patches/README.md for the gate and env knobs). Block 13's fused
MoE MMQ gate now covers RDNA4 + RDNA3_5 + RDNA3_0 (gfx1151 validated
2026-09-05, gfx1100 validated 2026-09-05 — see
Current state).
├── README.md # this file: overview + consumer workflow
├── AGENTS.md # working guide for LLM agents in this repo
├── MANIFESTS.md # apply order, per-block verification, validation history
├── BASELINE.md # fork point, patch provenance, drift policy
├── GREEDY-PURITY.md # purity rulebook: index, invariants, per-finding claims (read before shipping)
│ # narratives/evidence for the closed cases: archive/docs/GREEDY-PURITY-FINDINGS.md
├── WORKLOG.md # dated delivery records (newest first; README points here)
├── rdna-boosts-all.patch # convenience: the entire 16-patch net as ONE patch
├── patches/ # the delivery set: 0000-0015
│ └── README.md # apply instructions + block-12 env knobs + server config
├── scripts/
│ ├── apply-all.sh # the verified apply flow (git am; automatic -3 fallback on drift)
│ └── make-patches.sh # regenerates the set from the fork (~/llama.cpp)
├── benchmarks/ # benchy methodology + v1/v2 results + graphs (dated records)
├── prompts/ # versioned, hash-stable test prompts (sha256-recorded; never edited in place)
├── wiki/ # source for the GitHub wiki (Home, MTP & Adaptive MTP, Quick Reference); see wiki/README.md
├── wip/ # ACTIVE exploration docs / handoffs (currently only: nwarps/)
├── upstream/ # upstream-PR candidates (UPSTREAM-PR-*.md + .patch) + their index
└── archive/ # the rest: archive/work/ (closed experiments + the archived wip/ trees) + archive/docs/ (history)
History: the
baseline/<sha>branches,block/01-…11tags, and all dated validation records belong to the old pre-block-12 structure and live inarchive/docs/(see alsoarchive/work/for the closed experiments). Do not mix them with the currentpatches/files.
| patch | what |
|---|---|
0000 |
structural and architecture fixes — FA small-batch KV-split width invariance (issue #25) + Vulkan masked-V/freed-cell fixes (dead columns never read V). The base every later block applies on top of. |
0001 |
adaptive MTP draft depth (--draft-mtp-adaptive) |
0002 |
fused chunked gated-delta-net prefill kernel (bf16/WMMA, arch-segregated gfx12/gfx11) |
0003 |
BF16 KV cache + native-BF16 flash-attn (+ the HIP masked-V/freed-cell fixes since 2026-09-10) |
0004 |
RDNA4 WMMA flash-attn + Q6_K mmq prefill perf (WMMA path also runs on RDNA3.0/3.5, tuned head limits) |
0005 |
CPU bit-identical decode/verify batches |
0006 |
general system-operations bucket — the delivery's catch-all for changes that fit no other block: the FA instance build-time work, --fit under -sm tensor, the host-buffer input layer, the tiny-CPU-split single-thread fix and (r12) the op-offload H2D staging ring + tensor-split op-offload. Named for what it is since r12; the original host-buffer revert content is long gone (upstream reverted #24233 in #28604). Amended r13 (issue #52): the tiny-CPU-split heuristic now counts what a graph reads — not only its node outputs — so a CPU-offloaded FFN chunk keeps the full thread pool |
0007 |
meta device-wrapper skip |
0008 |
fused-core prefill kernels + GPU bit-identical results (needs blocks 03+04; amended 2026-09-07 with the mul_mat+add through-view shape guard, PR #15). Now folds the mmb (bf16-WMMA dequant weight GEMM) core, the RDNA4 fragment port / per-arch tuning, and the GDN/PLE conv1d + narrow-row RMS-norm prefill fusions (absorbed from the former archive/work/mmb-general campaign). |
0009 |
meta-buffer compute-container headroom |
0010 |
k-quant-boosts: Q4_K/Q5_K/Q6_K/Q8_0 mmvq VDR (+ q8_1 quantize-cache fusions; adds a dedicated RDNA3.5 mmvq table) |
0011 |
skip CUDA graphs for multi-token PRE-FILL (decode keeps graph replay) |
0012 |
hybrid HIP all-reduce — custom internal AR for the small-tensor decode path, per-size hybrid dispatch vs RCCL, RDNA4-only gate (bounded in-kernel spin since 2026-08-30 fix round; builds without RCCL) |
0013 |
fused MoE gate+up+GLU MMQ + mmvq short-K item-split — prefill fused expert MMQ (RDNA4 + RDNA3_5 + RDNA3_0, Q3_K/Q4_K/Q5_K/Q8_0/Q6_K, env opt-out GGML_CUDA_DISABLE_MOE_MMQ_FUSION) + decode item-split (rpb 2/4/8) merged with the upstream has_fusion mmvq path |
0014 |
qwen4exp / Qwen3.8-Flash-Next support — QSA sparse FA (default) + fused indexer top-k, HC_MIX/HC_COMBINE fused decode ops, managed lazy reader, MTP draft-head, WS4 hyperconn prefill fusions, QSA decode campaign + per-arch dense/QSA decode policy (promoted from beta/qwen4exp; see patches/README.md block-14 notes). The masked-V/freed-cell fixes it once carried now live in blocks 00 (Vulkan) and 03 (HIP). |
0015 |
attention-memory wins (block 15) — promoted 2026-09-12 from archive/work/block-15-campaign-wins/: V3 derived kq mask (LLAMA_KQ_MASK_DERIVED, on by default), V4 native q8_0 + V5 native bf16 K/V in the FA kernels (both behind GGML_CUDA_FA_KV_NATIVE, opt-in default 0), W1 QSA score-chain memory (GGML_QSA_SCORE_MEM), W2 derived QSA per-block bias + visibility (GGML_QSA_DERIVED_BIAS/GGML_QSA_DERIVED_VIS), W3 keys-only QSA indexer cache (LLAMA_QSA_KEYS_ONLY), W4 ggml-alloc unused-view release (no gate; A/B revert in archive/work/block-15-campaign-wins/ab/). ~3.4 GiB/GPU + ~1.2 GiB host saved on qwen4exp, ~800 MiB/GPU + ~800 MiB host on dense models, at ~1.3 % prefill / ~0.3 % decode. Now also folds the qwen4exp/QSA/HC/indexer campaign (qsa3 packed-block WMMA attention, fused indexer top-k + prefill score fusions, HC16 native-BF16 producers, hc_gate_mix, sparse MTP-draft attention, the sparse-QSA/derived-indexer defaults, and the host-buffer/CPU/meta fixes) — the former archive/work/mmb-general work. |
Block 15 (attention-memory wins) is part of the delivery since 2026-09-12 (
patches/0015, promoted fromarchive/work/block-15-campaign-wins/; a fresh set is now 16 patches, blocks 00-15).
Greedy-purity note (read before shipping): on the K-split decode paths, block 10 (
0010) is the only patch that changes decode numerics on ANY architecture — its VDR kernels reorder the fp32 reduction. Compute outputs are not bit-identical to a build without it (max logit diff 0.184 vs 0.203 for flash-attn on/off; greedy streams are deterministic within a build but can flip across configs). This is a different rounding path, not a correctness change. If you require 100% greedy purity across builds, do not install0010-…k-quant-boosts…patch— it is one line to drop fromscripts/apply-all.sh. Full discussion:GREEDY-PURITY.md. Block-13 caveat (2026-09-02): block 13 rewrites the small-batch mmvq decode kernel and is a second decode-numerics source on the rows that run it (short-K K<4096 ncols==1 rows, MoE projections; ncols 2..8 and long-K rows were restored to the pre-block-13 K-split kernel by the 2026-09-02 fix). Excluding block 10 no longer reproduces stock bits exactly on those rows — see GREEDY-PURITY.md §9.
# 1. fresh clone of llama.cpp, at the fork point recorded in release.json
BASE=$(jq -r .base release.json) # from this repo
FORK=https://github.com/ggml-org/llama.cpp
git clone $FORK && cd llama.cpp
git checkout "$BASE"
# 2. apply the set (automated; strict 16/16 git am on the recorded base)
bash <path-to-this-repo>/scripts/apply-all.sh .
# = git am patches/0000…0015 (one commit per block on a fresh `rdna-boosts` branch)
# 3. build + verify (trim -DGPU_TARGETS to your GPU arch for a faster build)
cmake -B build -DGGML_HIP=ON -DGGML_HIP_RCCL=1 -DGPU_TARGETS="gfx1100;gfx1151;gfx1201" -DCMAKE_BUILD_TYPE=Release
cmake --build build -j
# coherence gate (same-seed output must match a known-good build):
./build/bin/llama-cli -m <model> -ngl 99 -sm tensor -mg 0 -p "The capital of France is" \
-n 20 --seed 42 --temp 0 --no-display-prompt --single-turn
Speed up rebuilds with ccache. The set's flash-attention template instances are the build's critical path, and their native-KV loader arms are deliberately force-inlined — the optimiser's cross-inlining is what makes them fast at runtime and slow to compile. With
ccacheon PATH, a wiped rebuild of unchanged sources is a full cache hit: measured 282 s -> 4.2 s on a 16-core gfx1201 box, 321.8 -> 5.2 s on gfx1151 and 383.9 -> 4.8 s on gfx1100 (657/657 compile steps hit on each). Add-DCMAKE_HIP_COMPILER_LAUNCHER=ccache -DCMAKE_C_COMPILER_LAUNCHER=ccache -DCMAKE_CXX_COMPILER_LAUNCHER=ccache(the launcher form works with ROCm clang HIP device compilation; ccache 4.12.3 tested, on CMake 4.3). ccache replays the compiler's own objects, so the cached build is the same code — verified with same-seed greedy text (identical hash on every host before and after enabling it),llama-bench(within noise) andtest-backend-ops. Any header change (e.g.fattn-mma-f16.cuh) invalidates its dependents, i.e. the whole FA group.Do not use
git applyon the concatenated 1-11 series — it silently drops hunks (30 files / 2483 lines vs the correct 35 / 6094, verified 2026-08-29).git am(orscripts/apply-all.sh) is the required flow.
git am patches/000[1-9]-*.patch patches/001[0-5]-*.patch # blocks 01-15
git add -A && git commit -m "rdna-boosts: block 15: campaign memory wins"On the beta-integration branch the mmb (bf16-WMMA dequant weight GEMM) / QSA / indexer
campaign — the former 28-patch archive/work/mmb-general/ set — is folded directly into the 16 delivery
blocks, so the normal workflow above is all there is to apply. There is no separate beta layer
any more:
- block 08 absorbs the
mmbcore, the RDNA4 fragment port and per-arch tuning, and the GDN/PLE conv1d + narrow-row RMS-norm prefill fusions; - block 06 (the catch-all system-operations bucket) absorbs the host-buffer input layer and the tiny-CPU-split single-thread fix;
- blocks 13/14 absorb the
mmbfusion stand-downs and the extended MMVQ routed band; - block 15 absorbs
qsa3, the fused indexer top-k + prefill score fusions, HC16,hc_gate_mix, sparse MTP-draft attention, the sparse-QSA/derived-indexer defaults, and the meta/CPU backend fixes.
Applying the 16 patches to 84e76d8a2 therefore reproduces the full campaign tree (r8's
24bb0f5acb3e866abd4cad8c0de1bad45a20cb47, plus r9's issue-#47 MMA-FA typed-store fix -> current
a3dc4bbb680bf9dd8bcb5949ec833dec2a892aeb) in one pass:
git clone https://github.com/ggml-org/llama.cpp && cd llama.cpp
git checkout 84e76d8a2
bash <path-to-this-repo>/scripts/apply-all.sh . # 16/16 strict, tree a3dc4bbb…The per-patch fold mapping and validation record are in
archive/work/beta-integration/integration.md. Much of the MMB
kernel work is heavily adapted from pwilkin's
strix-halo fork, with thanks.
Historical record only.
archive/work/mmb-general/(the 28 patch files, themmb-general.patch,BETA-TESTING.md, the gfx1201/gfx1100 records) is kept as the campaign's verification record; its patches are no longer applied separately and theapply-beta.shhelper has been removed (the delivery itself now contains the campaign). Its gfx1151 beta-window re-validation was GREEN (2026-09-25 — the four gates + the recurrent rollback; seeBETA-TESTING.md§8), which is what the fold relies on.
The best general-purpose speculative-decoding configuration measured on this delivery combines the
adaptive MTP controller with the draftless ngram-mod speculator:
--spec-type draft-mtp-adaptive,ngram-mod \
--spec-ngram-mod-n-match 45 \
--spec-draft-n-max 9 --spec-draft-n-start 9ngram-mod supplies the long verbatim-recall drafts the MTP head cannot match, while the adaptive
controller keeps the MTP depth right for everything else. Measured against plain draft-mtp-adaptive
on the Q8_0 2-GPU reference cell (-n 3000, a 4-prompts-per-axis corpus): recall +67.5 %,
code +0.8 %, prose +0.5 %, reasoning −1.9 %, overall +13.6 %; on the dense 1-card cell it is
code/prose-neutral with the same recall win. The small reasoning cost is the price of the deeper
n-max 9; a workload with no verbatim recall is marginally better served by plain
draft-mtp-adaptive.
Guidance:
- Cap.
9is the general-purpose pick and the optimum on a single card (a cap of6costs 11-14 % on code there). On a multi-card tensor split6-7is ~1 % better. A lower cap saves only a small verify-batch scratch, not the model/KV memory. n_match. Keep it>= 40and an integer multiple of the cap —9/45,8/48,6/42. Too short (nm24) makes ngram fire on incidental code repeats and lose code throughput; a non-multiple (e.g.nm42at cap 12) degrades acceptance.- The combo changes the draft strategy, so it is opt-in — the controller default stays plain
--spec-type draft-mtp-adaptive.
Full derivation (the 4×4 corpus, all four cells, and the rejected controller alternatives):
archive/work/mtp-journey-2026-09-17/SUMMARY.md (narrative in its
README.md); the dated controller records are in
benchmarks/, newest 2026-09-15-adaptive-mtp-tuning.md.
On by default, deliberately. The delivery derives the attention mask inside the FA kernel from
compact per-cell state instead of materialising the n_kv x n_q f16 mask. That removes
n_ubatch x n_ctx x 2 bytes of compute-buffer VRAM plus the same again on the host — measured on
a 9B at -c 98304 / ub 512: 184.0 -> 88.4 MiB device and 112.0 -> 16.4 MiB host. The saving
scales linearly with the ubatch, which is the point: a deep-context MoE or qwen4exp /
Qwen3.8-Flash-Next workload wants a large ubatch, and that is exactly the configuration where the
mask is biggest (~800 MiB/GPU at ub 2048 / 196k) and where the VRAM the feature frees is the
difference between fitting the context and not.
The cost is prefill only — decode is untouched, because the derived path only fires for batches
larger than 8 tokens (speculative verify keeps the packed mask, so n_max <= 7 stays bit-identical).
Measured PP512, mask on vs off, -r 3 ("+" = the mask helps):
| config | d0 | 32k | 64k | 98k |
|---|---|---|---|---|
| gfx1201 9B dense 1 GPU | — | +1.9 % | — | +3.5 % |
| gfx1201 27B 2 GPU tensor | −1.3 % | −0.4 % | +1.3 % | +1.7 % |
| gfx1201 27B 2 GPU layer | — | — | — | −1.6 % |
| gfx1201 27B 3 GPU tensor | −3.4 % | −0.6 % | — | +2.0 % |
| gfx1151 9B dense 1 GPU | +0.4 % | −0.2 % | −0.9 % | −1.8 % |
| gfx1151 35B-A3B MoE 1 GPU | −0.3 % | −0.3 % | −0.8 % | −1.6 % |
| gfx1100 9B dense 1 GPU | −0.4 % | −0.5 % | −0.4 % | −0.2 % |
The tensor-split shape is the one to understand: the packed mask grows with n_kv, so on a
tensor split at shallow depth the mask is a small loss (−1.3 % at d0, crossing zero near 48k) and
becomes a win by 64k+; on the maintainer's 3-GPU tensor serving setup it is a win at depth. gfx1100
and gfx1151 pay a depth-growing ~1–2 % (they did not recover as much from the r7 kernel fix as
gfx1201 — the iGPU shares host bandwidth and the 7900 XTX has more of its own). gfx1100 on a
dual-card -sm tensor split is the one cell we still cannot measure here (only a single 7900 XTX
is available); a community report on 2x RX 7900 XTX is pending.
Turning it off. LLAMA_KQ_MASK_DERIVED=0 restores the packed mask (upstream's behaviour).
Worth doing if you are on gfx1100/gfx1151 and want the last ~1–2 % of deep prefill, or on a
tensor split at shallow depth and prefill latency matters more than the VRAM. For a deep-context
MoE / qwen4exp workload the default is the right side of the trade.
Not a correctness knob: same-seed output is byte-identical either way (the derived mask produces the same values; only the memory layout and prefill cost differ).
It works on both prefill kernels (r9). The derived mask is implemented by the MMA and the tile flash-attention kernels, so it is no longer tied to the chooser picking MMA: a head above the per-arch WMMA cap (RDNA4 576, RDNA3_5 320, RDNA3_0 256) used to lose the mask entirely, and that is every Gemma4 (head 512) on gfx1100/gfx1151 — the two arches that live on the tile kernel. There the mask is a win, not a tax (PP512, mask on vs off):
| config (tile kernel, natural selection) | cell | delta |
|---|---|---|
| gfx1100 Gemma4 12B, q8_0 KV | @ 16k / @ 32k | +1.2 % / +0.9 % |
| gfx1151 Gemma4 12B, q8_0 KV | @ 16k / @ 32k | +1.6 % / +0.6 % |
| gfx1201 Gemma4 E4B (tile forced) | @ d0 | +2.3 % |
and decode pays nothing for it. Decode and the spec verify batch always take the tile kernel (the
chooser's WMMA branch requires ne[1] > 8), so the derived branch there is a cost on every arch; it is
hoisted out of the KV loop so the packed path's code generation is unchanged. Measured r8 vs r9 at
that kernel: tg128 deltas of +0.01 % (gfx1201 9B @ d16384), +0.02 % (gfx1100 9B @ d16384),
and flat on gfx1151 — the earlier per-iteration form cost −0.5..−0.8 % at depth before the hoist.
The vec kernel has no derived arm, but it is decode/verify-only (n_tps <= 2) while the derived
form only exists for prefill-shaped batches (kq_mask_derivable() rejects n_tokens <= 8), so it
cannot be selected for one. If the launch log does print derived kq mask flash attention not supported, set to disabled on a CUDA/HIP backend, the FA node did not reach the GPU at all — check
which kernel serves that head (the log adds a note pointing there).
One performance caveat with nothing to do with this knob: forcing the tile kernel for a head it
would not normally serve (a stale GGML_CUDA_FA_WMMA_256=0, a fixed env in the September qwen4exp
gates, is the usual cause) makes a head-256 model on gfx1201 ~3x slower at deep prefill (9B,
-d 98304: 2104 -> 710 t/s). That env is worth removing regardless of the derived mask. One
exception worth knowing: qwen4exp / Qwen3.8-Flash-Next gets its deep-context mask elision from the
QSA path's own derived visibility (GGML_QSA_DERIVED_VIS, the code's "-800 MiB win"), which is
independent of this knob; LLAMA_KQ_MASK_DERIVED only serves that model's dense shortcut
(n_kv <= 2051), where the mask is tiny.
Full matrix, raw CSVs and the A/B harness: archive/work/kq-mask-derived-ab/; the
2026-09-19 block-15 (r7) amendment in patches/README.md.
The patches are static against the fork point in release.json.base. When upstream
drifts and hunks no longer apply, re-base the block commits (the fork checkout carries
them), regenerate the whole set with scripts/make-patches.sh, then refresh
release.json (scripts/make-release.sh --base … --tip … --tree …) and update the
current-state headers. The old baseline/<sha>-branch-per-upstream-range workflow
was retired when the delivery moved to the flat 16-patch set on main.
Some blocks are candidates for upstream contribution to
ggml-org/llama.cpp; others are
expected to stay fork-local. Block 12's internal all-reduce is gated to
RDNA4 pending community verification on RDNA3 pairs. See MANIFESTS.md
for per-block verification and BASELINE.md for provenance.
- 16-patch set (block 00 + blocks 01-15) for llama.cpp at the fork point
84e76d8a2(upstream master "metal : fix graph capture and handle empty graphs", 2026-09-24 re-base). - Canonical 16-block chain on
main: tip20b0efc5b273b26f6892012edb07d81e08b44d30, net treedc7ce12a6af627b0f140b9743e62bc0204f11b10(r8 campaign tree + the issue-#47 store fix + the r10 mask skip + the r11rpbmis-launch fix + the r12 staging ring + the r13 tiny-graph fix + the r14 derived-mask device-window fix + the r15 host-expert pinning fix + the r16 host-resident-expert prefill fast path + the r17 decode regression fix); releasev16-84e76d8a2-r17. - Host-resident MoE decode is restored (block 06, r17, 2026-09-27): r13's tiny-CPU-graph heuristic
counted the bytes a graph reads, which is right for a CPU-offloaded dense
MUL_MATFFN chunk but also countedMUL_MAT_ID's src0 — the whole expert weight table (144 MiB) — so the offloaded decode MoE graph (~120 one-tokenMUL_MAT_IDgraphs per pass under-ncmoe) stopped being classified "tiny" and ran multi-threaded, where the per-graph thread-pool re-arm dominates the work.MUL_MAT_IDsrc0 is now exempt exactly likeGET_ROWSsrc0: Qwen3.6-35B-A3B Q8_0-ncmoe 99tg6413.8 → 24.4 t/s (the r12 baseline), pageable 13.5 → 22.7, with the ±1.25 variance gone; the dense-MUL_MATissue-#52 path (8.7 → 1.7 t/s) is untouched and same-seed output is bit-identical. - Host-resident MoE experts (
-ncmoe) under-sm tensorare now a first-class prefill configuration (block 15, r16, 2026-09-27): the op-offload H2D staging ring is on by default (GGML_SCHED_STAGE=0opts out), the metastage_inputgained a split branch that gathers each device's slice into its ring slot on the copy stream (its slice-sum guard had compared against the whole-tensorsizeinstead ofchunk_size_full, so-sm tensoroffload had silently fallen back to the slow splice), the compacted strided splice uses a pinned gather + queued 1-D H2D (GGML_CUDA_SPLICE_GATHER=0reverts) instead of the pageablehipMemcpy2DAsync, and split expert copies are on (GGML_META_SPLIT_COPY=0reverts). All self-select from-sm/-ncmoe, so a stockllama-server … -sm tensor -ncmoe Nneeds no env vars. On Qwen3.6-35B-A3B Q4_K_M (gfx1201 x1/x2, pp8192) the default 2-GPU-sm tensor -ncmoebeats upstream84e76d8a2at every offload level — +91 % at-ncmoe 0rising to +148 % at-ncmoe 40(all experts host) against upstream's only 2-GPU option (-sm layer) — andtensorbeatslayerby +21 % → +33 %. Bit-identical output;MUL_MAT_ID/FLASH_ATTN_EXTgreen. - Host-resident MoE experts now stay pinned (block 06, r15, 2026-09-27): with
-ncmoethe scheduler op-offloads the used experts every ubatch, butselect_weight_buft's "avoid using a host buffer when using mmap" downgrade sent those uploads through the pageable model mapping — which on ROCm blocks the host insidehipMemcpyAsync(so the two cards' DMAs cannot overlap) and makes the meta backend's 2-D spliced upload fault inhipMemcpy2DAsync. The downgrade is now skipped forMUL_MAT_IDweights (default on,LLAMA_MMAP_HOST_EXPERTS=0restores it), which is +83 % on-sm tensor -ncmoe 99pp8192 (2794 -> 5104 t/s on 2x R9700) and bit-identical output. Cost: the expert set is pinned, non-swappable RAM. Same finding, independently, in GenerelSchwerz'smoe-cachefork. - The derived kq-mask inputs are no longer read on the host (block 15, r14, 2026-09-27, issue #53):
the r10 fully-masked KV-group skip built its batch-wide bitmap by dereferencing the derived
tok_lo/tok_hifrom the CPU inlaunch_fattn, but the backend scheduler copies those host graph inputs to the compute backend, so the launcher sees device copies — a host read of device memory is0xC0000005inggml-hip.dllon the first prefill ubatch wherever the allocation is not CPU-mapped (Windows/WDDM; Linux masks it) and a race against the in-flight copy elsewhere. The window is now reduced cooperatively insideflash_attn_kq_derived_blocks; the bitmap is unchanged, so the skip is still exact. The same fix replaces the directt->datawrites in thederived/mask_hole/FLASH_ATTN_QSAtest-backend-opsinitializers withggml_backend_tensor_set(the firstderived=1case segfaulted on Windows before any derived case ran).FLASH_ATTN_EXT6354/6354 and same-seed text identical across{skip on, skip=0, derived=0}; interleaved prefill A/B within noise. SeeWORKLOG.md(2026-09-27 r14) andpatches/README.md. - The tiny-CPU-graph single-thread heuristic no longer serializes CPU-offloaded FFN decode (block 06,
r13, 2026-09-27, issue #52): the heuristic added in r8 summed only the nodes' output activations
when deciding "tiny", so a CPU-offloaded FFN chunk (four
MUL_MATnodes with ~16 KiB outputs that read tens of MiB of weights each token) ran on a single thread — 8.66 -> 1.74 t/s on the reporter's box, and 3.22 -> ~4.9 t/s on the gfx1201 + 9950X3D reproduction. It now counts each node's non-view input tensors too, withGET_ROWSsrc0exempt (the embedding table is read only where gathered), so the host-resident-embedding spec-decode case the heuristic exists for keeps its single thread. Greedy same-seed output is byte-identical with the heuristic on or off. SeeWORKLOG.md(2026-09-27 r13) andpatches/README.md. - The op-offload H2D staging ring is delivered (block 06, r12, 2026-09-27, issue #50 / PR #51 by
@briansp2020, whose redirect design the merge adopts): a prefill's host-resident MoE expert upload
is issued on a per-device copy stream into a bounded slot ring and overlapped with the previous split's
compute, with a link-calibrated width gate and a fallback that disables staging rather than degrading
(measured
pp8192+81 % on the x4 box, +31-81 % on his 55 GB/s box; same-seed text, MTP,W=1..8and perplexity all byte-identical with staging on or off). The same block fixes the-sm tensorop-offload path — the meta device declared nooffload_op, so-ncmoehad been executing the whole MoE on the CPU (523 -> 1823 -> 2742 t/s atpp8192/ub8192) — and completes 15 backends'ggml_backend_iinitializer lists, a latent bug that had NULLedgraph_optimizefor metal/vulkan/hexagon/virtgpu. Block 06 is renamed togeneral system-operations bucket. SeeWORKLOG.md(2026-09-27 r12) andpatches/README.md. - The MMA FA prefill K/V store is typed again on the non-swizzled (AMD) path (block 15, r9,
2026-09-26, issue #47): upstream
1884824fd's swizzle refactor left the generic loader storing through(char *) tile_KV + swizzle_bytes<…>, which is address-identical but drops thehalf2alignment, so HIP split the 16-byte shared store. Restoring the typed store underif constexpr (!swz)recovers +2-7 % onhsk=256prefill shapes and +4.4 % on the reporter's 27B UD-Q4_K_XL q8_0 pp4096 @ d40000 (872.6 -> 911.3 t/s; reporter 885.5 -> 922.0 = +4.1 %), decode flat, output byte-identical. Seepatches/README.md(2026-09-26 block-15 r9) andWORKLOG.md. - The
mmb/QSA/indexer campaign is folded into the delivery (2026-09-25, releaser8): the former 28-patch opt-inarchive/work/mmb-general/set is now part of the 16 block patches — themmb(bf16-WMMA dequant weight GEMM) core, the RDNA4 fragment port / per-arch tuning and the GDN/PLE/RMS prefill fusions in block 08; the catch-all system-operations fixes in block 06;qsa3, the fused indexer, HC16,hc_gate_mix, sparse MTP-draft and the MMVQ band in block 15 (with block 14's pair stand-down and block 13's GLU stand-down). Strictgit am16/16 reproduces the full campaign tree24bb0f5acb…and the gfx1201 build is clean.archive/work/mmb-general/is kept as the historical verification record and theapply-beta.shhelper has been removed. Seearchive/work/beta-integration/integration.md. - The RDNA4 GQA-6 decode/verify FA band covers f16 (and, through its native arm, bf16) too
(block 15, r5, 2026-09-25, issue #45 follow-up, reported by
@DanoPTT): the
head-256 GQA-6
n_q <= 8band no longer pays the tile kernel's 3x K/V re-fetch/dequantization; the whole band runs the WMMA kernel with the GQA group folded into one block (ncols2 = 8) and the KV split round-robin over a fixedP(independent ofn_qand the KV length), so decode and every verify width reduce identically. It started quantized-only (r4, reported by @overdoingism) and r5 extends it to the 2-byte types with a per-element-size config: native-quantized keepsncols1 = 4/P = nsm, the 2-byte types takencols1 = 2/P = max(2, 3*nsm/4). f16 kv 102400 verify widths 2.1-3.2x faster; 27Bdraft-mtp n3at ~30k +13 % (f16) / +14 % (bf16 native); plain f16 decode -2.5..-4.2 %. r6 (2026-09-25) flips the bf16 native K/V arm to default ON (GGML_CUDA_FA_KV_NATIVEunset now enables it,=0disables every native arm), since r5 made native bf16 the path to the band and the incoming beta prefill boosts outweigh its small prefill cost. Default on (GGML_HIP_FA_BAND_WMMA=0opts out); prefill untouched. Seepatches/README.md(2026-09-25 block-15 r4/r5/r6) andWORKLOG.md. - The qwen4exp CPU
hc_combinereference is correct (block 14, r3, 2026-09-25, issue #44):ggml_compute_forward_hc_combine_f32readblock_outwith at*ne[1]row stride andinjectwitht*hc, but the model hands both over as multi-token tensors whose ownnb[1]differs (block_outis a contiguous[n_embd, nt],injecta view into the mix output with row striden_embd+hc). Every fused multi-token ubatch therefore read the wrong rows and a CPU-resident qwen4exp decoder layer emitted EOS as its first generated token; nt == 1 was accidentally correct. The reference now uses each tensor's ownnb[1](0 for a broadcastne[1] == 1), mirroring the CUDA kernel, and is bit-identical at nt == 1. Validated with an op-level CPU-vs-HIP oracle that fails 7/8 multi-token cases pre-fix and passes all 8 post-fix (gfx1100). - The wide-VDR MoE expert path is RDNA4/RDNA3_0-only (block 10, r2, 2026-09-25): the
VDR_Q4_K/Q5_K/Q6_K_Q8_1_MMVQ_MOEentry points were unconditional while only Q8_0 was arch-gated, so RDNA3_5 (gfx115x) ran the Q4_K/Q6_K experts — the Q4_K_M expert types — with the wide chunk the block-10 comment reserved for RDNA4/RDNA3_0.get_vec_dot_q_cuda()/get_vdr_mmvq()now ignoremoeon every other target in one place; base-16 MoEdraft-mtp n30.73967 → 0.76484, 87.5 → 89.6 t/s. - Shared-NextN MTP heads are usable (block 00, r13, 2026-09-22): a head with
nextn_shared_target_tensors(notoken_embd/outputof its own, e.g. the qwen4expmtp-…-shared-Q8_0.ggufsidecar) died every draft round on the M-RoPEX < Ycheck because the MTP driver inferred KV sharing fromctx_otheralone.is_mem_sharedis now gated on thegemma4-assistantarch; it is an upstream bug (04eb4c446, #23398) folded into the block-00 base. --fitworks under-sm tensor(block 6, r12, promoted from the formerbeta/tensor-fit-fix/, nowarchive/work/tensor-fit-fix/): upstream threwnot implemented for SPLIT_MODE_TENSORand swallowed it, so the default-on--fitwas a silent no-op under tensor split. The Meta device's accessors are now exposed andcommon/fit.cpphas a dedicated tensor path (per-device targets from--fit-target, a proportional split or an honoured-ts, then auto-n_ctxreduction and an-nglbinary search); an explicit-cis never overridden. Re-validated on r11 before promotion (fit decisions, 7 end-to-end loads with zero out-of-memory and zero compute-buffer growth, byte-identical same-seed gate).- The compute reserve accounts for the reachable (packed) kq mask (issue #42, block 15,
2026-09-20): V3's derived kq mask is a per-batch optimization, so a 2-D M-RoPE image/audio batch or a
multi-sequence batch allocates the packed mask (
n_kv*n_tokens*2bytes), which the reserve — measured with the derived form on — did not contain. At depth that mask is hundreds of MiB, so a deep-context image batch grew the compute buffer mid-run; under the default--fit-target 256that growth failed (cudaMalloc failed: out of memory,failed to process mtmd chunk) and the next request asserted.sched_reserve()now measures with the packed mask when such a batch is reachable (the newkq_mask_packed_reachable(): M-RoPE orn_seq_max > 1), so--fitcounts it exactly where it can happen. Same-seed output is byte-identical and throughput is unchanged; the reporter's M-RoPE model pays 8960 tokens / -4.4 % of fitted context, while a non-M-RoPE single-sequence model keeps V3's reserve untouched. A failed buffer allocation now also invalidates the allocator's layout instead of asserting on a later graph. - Block 11 replays HIP graphs for split-MoE decode again (issue #41, 2026-09-20): the pre-fill
test keyed off
nodes[0]->ne[1], which isn_expert_used(10) on the expert tensor a one-token decode split starts with under-ncmoe, so every decode split was skipped as multi-token. A newggml_cuda_graph_is_multi_token()reads the real token count fromMUL_MAT_ID'sne[2]/ a weightMUL_MAT'ssrc1->ne[1](0 -> 50 warmups / 0 -> 687 replays,tg10.6 -> 12.8 t/s on Qwen3.8-Flash-Next UD-Q4_K_XL, output bit-identical), and on HIP the exec is now destroyed/re-instantiated instead of updated, avoiding the ROCm <= 10.0hipGraphExecUpdateleak (GGML_HIP_GRAPH_FORCE_UPDATE=1opt-out). - FA instance build-time fix (blocks 06/13/15, 2026-09-18): the MMA instances are generated per
(ncols1, ncols2, head size)and the head-512 ones are listed first in the backend source order, the tile instances per(head size, KV type), and the fused-gate MMQ instances moved out ofmmq.cu. Cleanggml-hip -j16323.4 -> 236.0 s (-27 %), identical instantiations and symbols, no runtime change; the order, not the split, is what delivers it. - gfx1100 (RDNA3_0) WMMA FA is capped at head 256 (r5, block 04, issue #30): the 2026-09-14
RDNA4 #28102 config transfer shipped RDNA4-tuned rows and a lifted head cap to gfx1100, so head
512 took WMMA where stock takes tile and lost up to 23 % of deep prefill (gemma-4-26B-A4B
pp2048 @ d98304q8_0 661 -> 773 t/s, bf16 656 -> 851); head 256 keeps WMMA, a +44-52 % deep-prefill win. RDNA4 (576) / RDNA3_5 (320) are unchanged. - gfx1100 (RDNA3_0) tensor split keeps the stock AMD FA
ncols2rule (block 04, issue #30): the 2026-09-14 split-aware hint (wider genericncols2for tensor-split attention) was RDNA4-tuned and cost RDNA3_0 deep prefill (pp100K667.5 -> 779.4 t/s on 2× RX 7900 XTX, stock 805.0; decode unchanged). A single gfx1100 card is unaffected (it already took the AMD rule). --fitno longer SIGSEGVs with--spec-type draft-mtp-adaptiveand a minimal per-tier MTP head (issue #38; block 01, one line incommon/common.cpp).- A clean HIP build no longer prints the ~10k FA "loop not unrolled" warnings
(
-Wno-pass-failed, block 15; no codegen change). - Patches
patches/0000-…0015-…apply with strict 16/16git am(no 3-way fallback, whitespace-clean) viascripts/apply-all.sh.scripts/validate-set.shre-checks the artifact hashes, the strict apply and the applied tree againstrelease.json. hybridis the default all-reduce;GGML_CUDA_ALLREDUCE=ceselects the opt-in copy-engine (SDMA) 2-GPU mode and=ncclforces RCCL.- Greedy purity: plain decode ==
draft-mtpverify for--spec-draft-n-max <= 7across the supported KV types. Depths 8..15 are allowed with a visible notice (a verify wider than 8 rows switches kernel family);> 15is clamped (the recurrent rollback snapshot bound). - Last full
test-backend-opson this cut (gfx1201): 18083/18083, withFLASH_ATTN_EXT5952/5952 andFLASH_ATTN_QSA22/22.
The dated record of every change (re-bases, block amendments, issue fixes,
measurements) is WORKLOG.md, newest first. Per-block notes, env knobs
and server configuration live in patches/README.md; apply order
and the verification contract in MANIFESTS.md; fork point and drift
policy in BASELINE.md; the purity rulebook in
GREEDY-PURITY.md.
This work is becoming a community effort and I'd like to offer special thanks to the following users for the assistance in finding issues and offering solutions!
- https://github.com/1337hero
- https://github.com/bakon11
- https://github.com/briansp2020 (block-13 moe_weighted_reduction float4 remainder fix + block-14 MUL_MAT_ID pair-fusion layout gate, issues #19 and #18)
- https://github.com/eoprede
- https://github.com/overdoingism (issue #45: the RDNA4 head-256 GQA-6 decode/verify flash-attention band, reported with the diagnosis, op-level data, the round-robin KV split idea and a working opt-in patch; the r4 block-15 band is built on that submission)
- https://github.com/pwilkin (the
strix-halofork at https://github.com/pwilkin/llama.cpp/commits/strix-halo/, heavily adapted for the MMB bf16-WMMA dequant-weight GEMM work, now folded into the delivery — formerlyarchive/work/mmb-general/) - https://github.com/tungel
- https://github.com/DanoPTT (block-08 mul_mat+add through-view shape guard, PR #15; and issue #45 follow-up: the f16 verify-width diagnosis / the f16 + bf16 band coverage folded into block 15 in r5/r6, measured on their R9700)
I, and everyone else who benefits from this work, really appreciate you!
While most of the work in this repository are original works of my own, there are some significant portions, most notably around the prefill tuning, inspired by the excellent work performed by the community of: https://github.com/halo-box/strix-llama.cpp
Thank you to all the maintainers of the Strix Halo Llama.cpp project
Of course none of this would be possible without the baseline that all of this rests on, and that is the huge community over at https://github.com/ggml-org/llama.cpp
Many thanks to the llama.cpp team
Same as llama.cpp (MIT).