Skip to content

About

A set of performant RDNA/HIP patches against llama.cpp focusing on high speed inferencing with numerics assurance

Topics

assurancehipllama-cppperformancerdna3rdna4rocm

Resources

Readme
MIT license
Activity

Stars

203 stars

Watchers

5 watching

Forks

17 forks
Report repository
BranchesTags

Repository files navigation

llama-cpp-rdna-boosts

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.

Everything this delivery adds is default-on — you should not need to set any environment variable. ENVIRONMENT.md documents every variable the set adds or repoints: its default, and whether it is a kill-switch (default-on, set to disable), an opt-in, a tuning value, or a diagnostic.

It ships as 12 patches (block 00 + blocks 01-11) for a clean llama.cpp checkout at the fork point 097cf33d9 (upstream master). Each block is one git am commit, and the series is linear (a block assumes its predecessors), so the set is applied whole. The current release is v12-097cf33d9-r102: the release identity (fork point, canonical tip/tree, block count and per-artifact sha256) is release.json, and the release note for it -- and for every earlier rN -- is the top of WORKLOG.md (newest first; it is the only dated record). This README describes the set, not the changelog.

git clone https://github.com/ggml-org/llama.cpp && cd llama.cpp
git checkout 097cf33d9
bash <path-to-this-repo>/scripts/apply-all.sh .   # creates branch rdna-boosts (12 commits)

# OPTIONAL: a beta candidate on top (NOT the release tree -- for beta testing only)
bash <path-to-this-repo>/scripts/apply-all.sh --include-beta .   # 12 blocks + beta/*/*.patch

Releases

Frozen deliveries are published as GitHub Releases and tagged in this repo (the tag is the release identity: v12-<fork-point>-r<N>, e.g. v12-097cf33d9-r93; the v16-* tags are the historical 16-block set. r1 is the re-base of the set onto that fork point and every later rN is a dated change recorded in WORKLOG.md, newest first).

Supported architectures

The set targets the RDNA3 / RDNA3.5 / RDNA4 GPU families, with RDNA2 as best-effort:

family arches example parts status
RDNA 3 gfx1100 RX 7900 XTX/XT, RX 7800 XT, ... supported
RDNA 3.5 gfx1150/gfx1151 Strix Point / Strix Halo APUs supported
RDNA 4 gfx1200/gfx1201 RX 9060 XT; RX 9070 / 9070 XT supported
RDNA 2 gfx1030 RX 6800/6800 XT/6900 XT, Radeon Pro V620 best-effort; externally validated (V620), not run in our fleet

Strix Point (gfx1150) is in the target set as of 2026-10-05. Images published before then (for example the server-rocm-10.0 build at 2dbece4) carry no gfx1150 code object and fail on the first kernel launch (ROCm error: device kernel image is invalid). For those, build from source with gfx1150 in the target list, or set HSA_OVERRIDE_GFX_VERSION=11.5.1 at runtime to select the gfx1151 code objects. See Strix Point (gfx1150) for the validation and both paths.

RDNA2 (gfx1030) is best-effort: not a supported target, but no longer unvalidated. It is compiled into the image so RDNA2 owners pick up the generic, runtime-dispatched boosts instead of a first-kernel ROCm error: invalid kernel file (issue #158), and the RDNA3/4-only paths fall back to the upstream non-WMMA kernels at runtime. As of r76 the set also folds @slavap's tuned RDNA2 tile-FA and k-quant work (PR #165: per-DKQ tile config, launch caps, WGP-aware occupancy, the kv=113 partial-batch load fix and the wide-kernel Q-from-global path; Navi21 MMVQ rows/warps + MMQ float scale). All of it is gated to RDNA2, so gfx1100/gfx1151/gfx1201 keep their previous numerics. External validation (TODO #66, closed): a Radeon Pro V620 (gfx1030) passed test-backend-ops (FLASH_ATTN_EXT 6389/6389, MUL_MAT 2444/2444) and generated coherently with -fa 1 in a 3-card mixed gfx1100+gfx1030 -sm tensor split down to 32K depth. There is still no RDNA2 device in the maintainer's fleet, so there are no performance claims and no purity gates for it — an RDNA2 result is informational, never a delivery gate — and -fa on can still hit upstream RDNA2 tile-kernel defects, where -fa off is the workaround. See RDNA2 (gfx1030).

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), so a multi-arch build such as GPU_TARGETS="gfx1030;gfx1100;gfx1151;gfx1201" yields one binary that picks the right path on whichever of these it runs on. The WMMA kernel bodies are compile-gated by RDNA3/RDNA4; on RDNA2 the non-WMMA arms compile and the runtime predicates keep the WMMA kernels from being launched. 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 WORKLOG.md).

Layout

├── README.md              # this file: overview + consumer workflow
├── AGENTS.md              # working guide for LLM agents in this repo
├── MANIFESTS.md           # apply order + the verification contract (release gates per block)
├── BASELINE.md            # fork point, what the patches are, how they are cut, drift policy
├── TODO.md                # forward-looking tracker: open items, limitations, parked ideas only
├── GREEDY-PURITY.md       # purity rulebook: invariants + one summary entry per finding (read before shipping)
│                          #   narratives/evidence: archive/docs/GREEDY-PURITY-FINDINGS.md (same § numbers)
├── WORKLOG.md             # dated delivery records (newest first; the only release history)
├── ENVIRONMENT.md         # every env var the set adds or repoints, with its default and class
├── CONTAINERS.md          # the ROCm container images + the off-host build watcher
├── COMMUNITY-CONFIG.md    # running a model bigger than VRAM (`-ncmoe` + the expert cache)
├── rdna-boosts-all.patch  # convenience: the entire 12-patch net as ONE patch
├── patches/               # the delivery set: 0000-0011 (one `git am` block each)
│   └── README.md          # apply instructions + block-12 env knobs + server config
├── scripts/               # apply / regenerate / validate / publish tooling; index: scripts/README.md
│   ├── apply-all.sh       # the verified apply flow (git am; automatic -3 fallback on drift; --include-beta)
│   ├── gate-prefill-logits.sh  # prefill-logit KLD release gate (issue #113)
│   ├── gate-per-block-build.sh # a clean `all` build per block commit
│   └── 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 only, never applied; index: wip/CAMPAIGNS.md
├── beta/                  # promotion-candidate patches pending a beta test (NOT part of the delivery;
│                          #   apply with `apply-all.sh --include-beta`; currently empty, see beta/README.md)
├── upstream/              # upstream-PR candidates: one sub-directory each (README.md + patch) + index
└── archive/               # the rest: archive/work/ (closed experiments + the archived wip/ trees) + archive/docs/ (history)

History: the baseline/<sha> branches, block/01-…11 tags, and all dated validation records belong to the old pre-block-12 structure and live in archive/docs/ (see also archive/work/ for the closed experiments). Do not mix them with the current patches/ files.

The 12 blocks (V2.0)

release.json (V2.0; the current release). Each row is one git am commit.

patch subject concern
0000 foundation — band-uniformity, Meta-wrapper, allocator fixes old-00 + old-05 + old-07 + old-09a + the multi-ubatch logits readback + the ggml-alloc view release; + the meta hot-path dispatcher restoration (r86)
0001 hybrid HIP all-reduce (RDNA4-gated) the old block 12; + no per-AR ev.ker and the opt-in P2P-large path (r52)
0002 adaptive MTP draft depth (--spec-type draft-mtp-adaptive) old-01; + opt-in MTP prompt-tail ingest and the last-row prompt copy (r52); + the embeddings-layer-input preset API the draft setup calls (r94 block-boundary repair)
0003 chunked gated-delta-net prefill (bf16/WMMA GDN) old-02; + the per-sequence rollback tracker folds in every ubatch and LLAMA_RS_ROLLBACK_DEBUG (r91)
0004 BF16 KV cache + native-BF16 flash attention old-03
0005 flash-attention band / native quantized K-V / derived kq mask old-04 FA + the old-15 FA work; + the 2-byte arm joins the 32-column band block and the head-512 MMA row is retuned (r92); + the q8_0 fetch-ahead band variant, default on (GGML_HIP_FA_BAND_FETCH_AHEAD=0), r98 PR #195; + one tile K/V type selection for the instantiation and the F16 scratch, and an RDNA2-gated checked tail, r99 PR #198
0006 MMB quantized prefill GEMM old-08 MMB
0007 k-quant boosts (mmvq VDR, Q6_K MMQ, Q1_0, grid pad) old-10 + old-04 MMQ + old-15 Q1_0/grid/PR57; + the RDNA3_0 per-shape mmvq nwarps (r93)
0008 fused kernels (conv1d/RMS/top-k MoE/verify fusions + fused MoE gate+up+GLU) old-08 fusions + old-13 fused MoE + old-15 PR63/114; + topk_moe sorted-key argmax (r52); + the RDNA3_0 per-shape mmvq nwarps (r93)
0009 host-resident MoE (iface, scheduler, Meta, movable-boundary slab, cache engine, prefill cache staging) old-06 MoE + old-09b + old-11 + old-13 engine + old-15 hooks/arena; + the arena-host staging engine and its residency snapshot (r83); + the slab drop reference (r85); + the slab transient absorption common half and the derived draft reserve (r86); + the complete device-iface fills and the pre-loaded draft-model handoff (r94 block-boundary repair); + the batched f32 CPY run on RDNA3_0 (r95, PR #193); + the restored copy callback, so the upstream --moe-cache-mib engine works (r99, TODO #56); + the split-state walk's per-thread scratch arena, so a graph re-allocation stops allocating one heap node per visited tensor (r100, issue #184); + the --fit draft reserve charged only to the KV layers a drafter owns (r101), and a loaded-but-unused draft model is not loaded or given a context (r101)
0010 qwen4exp (model, hyper-connections, lazy reader, QSA/indexer) old-14 + old-15 HC16/policy; + short-K MoE rows-per-block (r52); + the stage_from_cache hook (r83); + the slab narrow band (r85); + the slab transient absorption CUDA/slab half (r86); + the async graph-input ring and its r90 split-input snapshot fix (default-on); + the fixed-D sequential GDN snapshot tail (r91, GGML_CUDA_GDN_TAIL=0); + the work-boundary shrink invalidates captured graphs (r94, issue #187); + the verify-band Q8_1 fusions and the GDN conv verify path on RDNA3_0 (r95, PR #192 + #193); + gemma4 dense -sm tensor (the Meta split state accepts an empty-or-whole slice for per-row ops, the rotated-K mul_mat and flash-attention K/V), the expert cache sizing / multi-slice tables / segmented async upload, the iSWA q, v, k KV-builder order and the per-origin async-input ring (r101)
0011 campaign remainder (CPU wins, DFlash, FA build grouping, mmap reclaim) old-14 QSA + old-15 W1-W3 + old-06/15 mmap reclaim + old-16; + the expected --fit ctx_other probe failure logs as a warning, and the draft-context creation is gated on an enabled type that uses it (r101)

Phase-2 additionally moved each concern into a self-contained file: the ggml_cuda_try_fuse arms, the MMVQ vec-dot/VDR selectors + MoE kernel (mmvq-vecdot.cuh, mmvq-moe.cu), the FA native-KV/band predicates (fattn-rdna.cuh), and the four ggml-cuda.cu TUs (ggml-cuda-try-fuse.cu, ggml-cuda-slab.cu, ggml-cuda-moe-cache-hooks.cu, ggml-cuda-arena.cu; ggml-cuda.cu 9819 → 6370 lines).

The pre-V2.0 16-block notes (provenance and edit history of the old block numbering) are in archive/docs/pre-v2-block-notes.md.

Consumer workflow

# 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 12/12 git am on the recorded base)
bash <path-to-this-repo>/scripts/apply-all.sh .
#    = git am patches/0000…0011  (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 ccache on 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) and test-backend-ops. Any header change (e.g. fattn-mma-f16.cuh) invalidates its dependents, i.e. the whole FA group.

Do not use git apply on the concatenated 1-11 series — it silently drops hunks (30 files / 2483 lines vs the correct 35 / 6094, verified 2026-08-29). git am (or scripts/apply-all.sh) is the required flow.

Manual equivalent

git am patches/000[1-9]-*.patch patches/001[0-1]-*.patch   # blocks 01-11

The mmb campaign is in the delivery

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 12 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 mmb core, the RDNA4 fragment port and per-arch tuning, the GDN/PLE conv1d + narrow-row RMS-norm prefill fusions, the mmb fusion stand-downs and the extended MMVQ routed band;
  • block 09 absorbs the host-buffer input layer and the tiny-CPU-split single-thread fix (the system work of the old catch-all block 06);
  • block 10 absorbs the deferred mmb pair-fusion stand-down and the qwen4exp group (the old block 14);
  • block 11 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 (the old block 15 remainder).

Applying the 12 patches to the fork point in release.json therefore reproduces the full campaign tree in one pass:

git clone https://github.com/ggml-org/llama.cpp && cd llama.cpp
git checkout 097cf33d9          # the base recorded in release.json
bash <path-to-this-repo>/scripts/apply-all.sh .   # 12/12 strict, tree checked against release.json

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, the mmb-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 the apply-beta.sh helper 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; see BETA-TESTING.md §8), which is what the fold relies on.

Recommended configuration — adaptive MTP + ngram-mod

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 9

ngram-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.

Do not combine ngram-mod with host-resident MoE experts. When MoE experts are partially or fully host-resident (-ncmoe/-cmoe, i.e. the MoE expert cache or the pinned ROCm_Host path is active), adding ngram-mod to draft-mtp-adaptive roughly halves decode throughput. Dropping ngram-mod and keeping plain --spec-type draft-mtp-adaptive restores the expected decode rate. The adaptive controller is itself a good fallback for replayed/verbatim content: it drafts deeper when the content is predictable, so it handles replay quickly (not as fast as ngram-mod, but far ahead of a halved decode). On a host-expert config, plain --spec-type draft-mtp-adaptive is therefore the recommended pick. The cause of the interaction is not yet understood; it is tracked in wip/ngram-mod-host-experts/. This applies only to host-resident experts: on a fully device-resident model the combo above is still the best general pick.

Guidance:

  • Cap. 9 is the general-purpose pick and the optimum on a single card (a cap of 6 costs 11-14 % on code there). On a multi-card tensor split 6-7 is ~1 % better. A lower cap saves only a small verify-batch scratch, not the model/KV memory.
  • n_match. Keep it >= 40 and 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. nm42 at cap 12) degrades acceptance.
  • Host-resident experts. On a config with host-resident MoE experts (-ncmoe/-cmoe), skip ngram-mod entirely (see the caveat above): the combination roughly halves decode, and plain --spec-type draft-mtp-adaptive is the better pick.
  • 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.

VRAM vs prefill — the derived KQ mask (LLAMA_KQ_MASK_DERIVED)

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.

Cross-start determinism on ROCm (issue #67)

Issue #67 (from #58 item D, reported by @DanoPTT on Windows / ROCm 10 / gfx1201) tracked Q6_K greedy output that intermittently swapped one near-tie across fresh llama-server starts while staying deterministic within a start. It is closed as an external ROCm issue — nothing to fix in this patch set. Two separate defects were in play:

  • A real bit-transparency bug in the fused rope path, fixed in r25. The address-gated ROPE -> VIEW -> SET_ROWS fusion did not reproduce the unfused chain bit-for-bit: clang contracted the fused <float,__half> and unfused <float,float> template instantiations differently, and one element of the prefill K-cache write landed on opposite sides of an f16 rounding boundary. r25's #pragma clang fp contract(off) in rope.cu makes every instantiation round identically, verified at the logits level on Linux (W=1 hash 60e77916673db071 with the fusion on and off). See WORKLOG.md for the r25 record.

  • The residual per-start flip is hipBLASLt solution selection in ROCm, not ggml — and it reproduces on a stock upstream build with none of these patches. The F32 mul_mat for ssm_alpha/ssm_beta (5120 -> 48) at 512 prefill columns misses the mmvf/mmf kernels, goes to hipblasSgemm, and rocBLAS routes it to hipBLASLt (rocblaslt_matmul, T,N, m=48 n=512 k=5120). HIPBLASLT_LOG_MASK=160 shows a per-process pick between a bit-identical reference solution pair (140231/140232) and a deviating pair (140216/140217); it is random from process to process, so layer 0's prefill output — and with it the whole run — differs at startup. This is ROCm/rocm-libraries#12126.

Workaround: ROCBLAS_USE_HIPBLASLT=0. On the reporter's setup, 150 fresh starts per arm: default 5/150 deviate, ROCBLAS_USE_HIPBLASLT=0 0/150 (bit-identical to the reference class), with no measurable prefill/decode cost on the current build (the earlier text-hash runs measured higher flip rates, up to ~14 %). ROCBLAS_DEFAULT_ATOMICS_MODE=0 does not help. The selection is prefill-only (decode at n=1 uses mmvf), so decode throughput is unaffected.

Linux / ROCm 7.14 does not reproduce it — the reporter's box (Windows / ROCm 10) is where the per-process selection actually varies. Linux with ROCm 10.x does (2026-10-21, gfx1201 / Linux / ROCm 10.1.0, which predates the upstream fix: gemma4 26B-A4B F32 MoE router, m=128 n=375 k=2816 T,N, solution_index 140231 vs 140232, 3 logit classes in 18 fresh starts of one MTP config; 10/10 identical once ROCBLAS_USE_HIPBLASLT=0). The fix is rocm-libraries c50ea60130f (TheRock nightlies from 2026-10-06). Keep hipBLASLt enabled — disabling it costs up to -15 % decode; set ROCBLAS_USE_HIPBLASLT=0 for a run that must be reproducible. The r24 rope-fusion kill switches (GGML_CUDA_DISABLE_ROPE_SET_ROWS=1, GGML_CUDA_DISABLE_RMS_NORM_MUL_ROPE=1) and r25's fix remain shipped for bisection and for the real defect above. Full record: GREEDY-PURITY.md §44. Second call site (2026-10-21, closes TODO #75): on 35B-A3B the same lottery rides the GDN F32 ssm_alpha/ssm_beta GEMM (F32 2048x32 x N): 3 logit classes in 8 fresh starts on ROCm 10.1.0, 1 class in 8/8 with ROCBLAS_USE_HIPBLASLT=0, the class being the TensileLite spine hipBLASLt picked — and there the workaround is free (+1-3 % prefill, flat decode, ~70 MiB less peak VRAM). Removing the hipBLASLt dispatch rather than pinning it is now a campaign: wip/hipblaslt-bypass/.

Concurrent serving with --kv-unified is not deterministic (issue #134)

Mixing -np > 1 with --kv-unified does not give run-to-run reproducible greedy output: a sequence's attention result depends on the KV-tile indices of its cells modulo the per-launch attention KV-split period (pb x nbatch_fa), and under unified KV every new decode token takes the next cell from the shared head, so that residue follows the co-resident session's allocation/scheduling history. Identical concurrent repetitions can differ, and a concurrent run can differ from the same request run solo, at the rounding level (a greedy near-tie can flip). Use --no-kv-unified when strict reproducibility is required. This is inherent to upstream's unified-KV design — it reproduces on unmodified upstream 097cf33d9 — and not a defect introduced by this patch set; the analysis, reproducers and the exact rule are preserved (closed 2026-10-19) at archive/work/kv-cell-offset-nondeterminism/ (RESOLUTION.md). A fix requires period-granular per-sequence cell ownership in the unified pool — an upstream allocator redesign with no throughput win, so it is not implemented. Single-request and non-unified runs are unaffected by the unified-KV effect, and the delivery's op-level purity invariants (W = 1..8, plain == draft-mtp, native == staging) are scoped to a fixed batch and are unchanged.

Separate from the unified-KV case (2026-10-16, issue #151 follow-up). The reporter (R9700 / Windows) also sees a start-to-start greedy flip with HIP graphs on, and reports that --no-kv-unified does not remove it. We could not reproduce it on Linux/ROCm 7.14: fresh-process decode-logit probes (W=1..8) are bit-identical across starts, graphs on and off, including on the reporter's 27B model with q8_0 KV (archive/work/kv-cell-offset-nondeterminism/FINDINGS-2026-10-16-revalidation.md). It is most likely the known Windows-only per-start flip (the r25 rope -> view -> set_rows bit-transparency fix, and the external hipBLASLt solution-selection issue, workaround ROCBLAS_USE_HIPBLASLT=0), both prefill-driven and independent of graphs. Confirmed and closed (TODO #67, 2026-10-20): the reporter ran the bisection on decode logits — with ROCBLAS_USE_HIPBLASLT=0 all starts are bit-identical, and with it on the same logit classes appear with graphs on and off, so graphs play no part. See archive/work/kv-cell-offset-nondeterminism/FINDINGS-2026-10-20-issue151-reporter-confirmation.md.

When upstream master moves

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 patch set on main.

Upstreaming

Some blocks are candidates for upstream contribution to ggml-org/llama.cpp; others are expected to stay fork-local. The internal all-reduce (block 01 in V2.0 numbering) is gated to RDNA4 pending community verification on RDNA3 pairs. See MANIFESTS.md for the verification contract and BASELINE.md for provenance.

Community Acknowledgements

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!

I, and everyone else who benefits from this work, really appreciate you!

Inspirational Works

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

License

Same as llama.cpp (MIT).

About

A set of performant RDNA/HIP patches against llama.cpp focusing on high speed inferencing with numerics assurance

Topics

assurancehipllama-cppperformancerdna3rdna4rocm

Resources

Readme
MIT license
Activity

Stars

203 stars

Watchers

5 watching

Forks

17 forks
Report repository

Releases

Packages

Contributors

Languages