cuda: AMD RDNA4 Q1_0/Q2_0 — HIP-path vec_dots (+37%/+16% decode), opt-in quant dedup, opt-in hipBLASLt prefill routes

[The-Monk]

What this is

The AMD side of the Q1_0/Q2_0 story. This fork's HIP release bundles currently run Prism's binary/ternary formats through kernels tuned for NVIDIA; on RDNA4 hardware that leaves a lot on the table. This PR contributes three self-contained pieces from our RDNA4 fork (The-Monk/llama.cpp roc8, the kernel base of The Rock8), re-measured on this tree on 2× Radeon AI PRO R9700 (gfx1201, ROCm/TheRock 7.14).

All numbers below: Bonsai-27B (qwen35 hybrid), 1 GPU, llama-bench -r 5 (tg) / -r 3 (pp).

1. HIP-path vec_dot for Q2_0 (+37%) and Q1_0 (+16%), on by default

The existing Q2_0 vec_dot extracts symbols with dynamic-selector __byte_perm chains, which are built around NVIDIA's PRMT; on HIP they lower poorly and the decode GEMV pays for it end-to-end. This PR adds an AMD path using the same split this file already uses for unpack_q1_0_bytes: bit-spread the raw 2-bit codes into bytes and use the identity dot(s,u) = dot(c,u) − sum(u), applying the offset once via the q8_1 stored sum — plain shift/mask + dp4a, no per-code arithmetic. The CUDA path is untouched.

The second commit gives Q1_0 the same treatment with the binary form of the identity (s = 2c − 1 ⇒ dot(s,u) = 2·dot(c,u) − sum(u)), replacing the select-chain fallback in unpack_q1_0_bytes that materialized ±1 bytes on a VALU-bound kernel.

tg128, Bonsai-27B, 1× R9700	stock	this PR	Δ
Q2_0 (ternary)	34.52 ± 0.08	47.20 ± 0.20	+36.7%
Q1_0 (binary)	54.53 ± 0.36	63.03 ± 0.70	+15.6%

Perplexity is identical (wikitext-2, 20 chunks: 11.5603 ± 0.47269 both builds, Q2_0), and binary now decodes faster than ternary, as its byte ratio says it should.

2. Opt-in activation-quant dedup (GGML_HIP_DEDUP_MMVQ_QUANT), default OFF

Sibling mmvq matmuls that read the same activation tensor (wq/wk/wv, ffn_gate/ffn_up, wqkv/wqkv_gate) each launch a quantize_row_q8_1 that recomputes byte-identical output. With the flag set, the quantization is computed once per activation tensor and reused; the cache lives on the backend context, is keyed by tensor pointer, and is reset at every graph_compute, so entries cannot survive a graph rebuild. GGML_HIP_DEDUP_MMVQ_QUANT_BATCH additionally widens it to ne11>1 (spec-decode verify batches).

Measured here: Q2_0 48.20 (+2.1% on top of #1), Q1_0 64.93 (+3.0% on top of the new Q1_0 dot). On our fork the same lever is worth considerably more under MTP self-speculative decode, where the verify pass multiplies the redundancy — that's why the _BATCH variant exists. Lossless: the cached bytes are the bytes the skipped launch would have produced; PPL run with both flags on is byte-identical to stock.

3. Opt-in Q1_0/Q2_0 hipBLASLt prefill routes, default OFF, optional dependency

Large-M (prefill) matmuls routed through hipBLASLt int8 GEMM (dequantize once → cached int8 weights + per-channel scale; activations int8 per call), with a self-tuning per-shape algorithm cache — gfx1201 has no hipBLASLt cost model and the heuristic picks badly. Includes the weight-cache invalidation registry (hipblaslt_wcache): caches are keyed on device addresses, so the backend drops entries when it frees a buffer — without this, multi-model runs and server model swaps can silently serve stale weights from a reused allocation.

warm tune cache	pp1024	vs stock kernel
Q2_0	1453.14 ± 0.39	+37.5% (1056.63)
Q1_0	1470.94 ± 1.74	+17.8% (1248.33)

Knobs per format (<FMT> ∈ Q1_0,Q2_0): GGML_HIP_<FMT>_HIPBLASLT_PREFILL=1 (enable), _MTHRESH=<n> (default 384 — deliberately conservative because the crossover is model-shape-dependent; override per model), _TUNE_CACHE=<path>, _NOTUNE=1, _WCACHE_MB=<n>, _FP8=1 (dequantize to e4m3 instead of int8).

Build impact: find_package(hipblaslt QUIET) — optional. Not found (e.g. Windows HIP, where hipBLASLt doesn't exist) → routes compile as inert stubs and the release bundles are unaffected. CUDA builds compile the same stubs (__HIP_PLATFORM_AMD__-guarded).

What we tried and deliberately did NOT include

Our fork runs Q2_0 with nwarps=8 / rows_per_block=3 on RDNA4. Transplanting those constants onto this tree's kernel was a disaster (nwarps=8 → 8.2 t/s) to neutral (full sweep: rpb 2/3/4 × nwarps 2/4/8 all ≤ +0.7% over stock). Launch-geometry constants do not transfer between kernel implementations; this tree's geometry is already at its optimum, so no tuning changes are included.

Testing

• test-backend-ops test -o MUL_MAT: all pass, both with default env and with both dedup flags set.
• PPL byte-identical stock vs patched(+dedup+batch), wikitext-2 20 chunks.
• Hardware: 2× Radeon AI PRO R9700 (gfx1201), TheRock ROCm 7.14, -DCMAKE_HIP_ARCHITECTURES=gfx1201.
• Not tested here: nvcc compile of the new files (no NVIDIA box on hand). The non-HIP sides are plain-C++ stubs / the untouched original code paths; happy to fix any CI fallout promptly.

Reproduce:

cmake -B build -DGGML_HIP=ON -DCMAKE_HIP_ARCHITECTURES=gfx1201 -DCMAKE_BUILD_TYPE=Release
cmake --build build -j --target llama-bench test-backend-ops
./build/bin/llama-bench -m Ternary-Bonsai-27B-Q2_0.gguf -ngl 999 -p 0 -n 128 -r 5
GGML_HIP_DEDUP_MMVQ_QUANT=1 ./build/bin/llama-bench -m ... 
GGML_HIP_Q2_0_HIPBLASLT_PREFILL=1 GGML_HIP_Q2_0_HIPBLASLT_TUNE_CACHE=/tmp/t.bin \
  ./build/bin/llama-bench -m ... -p 1024 -n 0   # run twice; first run pays the tuning cost

Standing offer

The dual-R9700 box these numbers come from is available for validating HIP release candidates or future AMD-touching PRs — happy to run your release CI artifacts against real gfx1201 before they ship.

🤖 Generated with Claude Code

[khosravipasha]

@The-Monk Thanks for the chnages, will do a more careful review this week.
Impressive changes, good to have better support for AMD.

By the way Q1_0 is fully upstreamed to llama.cpp so migth be worth also trying to upstream it there, so check whetehr they have added better AMD support since we merged things.
Similar for Q2_0, we fully upstreamed it but only caveat is that there group size is 64 so a bit different from our group size 128. More info here: https://github.com/PrismML-Eng/Bonsai-demo#upstream-status-for-ternary

Happy to merge these here but also worst trying to upstream to main llama.cpp as well since most things are already there too (with the caveat of Q2_0 being slightly different there).

[Copilot]

Pull request overview

Adds RDNA4 optimizations for Q1_0/Q2_0 inference in the shared CUDA/HIP backend.

Changes:

• Optimizes HIP vector-dot decoding.
• Adds optional activation-quantization deduplication.
• Adds optional hipBLASLt prefill routes, tuning, and weight caches.

Policy note: The PR description declares Claude Code generation, conflicting with AGENTS.md:36-40, which prohibits AI-written PR descriptions.

Reviewed changes

Copilot reviewed 11 out of 11 changed files in this pull request and generated 8 comments.

Show a summary per file

File	Description
ggml/src/ggml-hip/CMakeLists.txt	Adds optional hipBLASLt discovery.
ggml/src/ggml-cuda/vecdotq.cuh	Adds HIP/MUSA low-bit vector-dot paths.
ggml/src/ggml-cuda/mul_mat_q2_0_hipblaslt.cuh	Declares the Q2_0 prefill route.
ggml/src/ggml-cuda/mul_mat_q2_0_hipblaslt.cu	Implements Q2_0 hipBLASLt prefill.
ggml/src/ggml-cuda/mul_mat_q1_0_hipblaslt.cuh	Declares the Q1_0 prefill route.
ggml/src/ggml-cuda/mul_mat_q1_0_hipblaslt.cu	Implements Q1_0 hipBLASLt prefill.
ggml/src/ggml-cuda/mmvq.cu	Adds activation-quantization deduplication.
ggml/src/ggml-cuda/hipblaslt_wcache.cuh	Declares cache invalidation APIs.
ggml/src/ggml-cuda/hipblaslt_wcache.cu	Implements invalidator registration.
ggml/src/ggml-cuda/ggml-cuda.cu	Integrates routes and cache lifecycle.
ggml/src/ggml-cuda/common.cuh	Adds FP8 conversion and cache state.

Suppressed comments (2)

ggml/src/ggml-cuda/mul_mat_q1_0_hipblaslt.cu:415

• The claim that all uses share one stream is false: the backend dispatches concurrent branches and devices independently. Both cache builders publish the entry immediately after enqueueing requantization, so another stream can hit and read it before initialization completes. Publish a completion event and wait on hits, or synchronize before insertion as the existing Hopper cache does at mmq-hopper-q1.cu:149-153.

// Returns cached int8 weight (building it on first miss if within budget), or
// nullptr -> caller must requant on-the-fly. Build + all uses share the stream,
// so the one-time requant is correctly ordered before any GEMM that reads it.

ggml/src/ggml-cuda/mul_mat_q2_0_hipblaslt.cu:425

• The cache entry is published immediately after requantization is enqueued. A concurrent branch or request on another stream can hit this entry and launch GEMM before the converted weights are initialized. Publish a completion event and wait on hits, or synchronize before insertion as the existing Hopper cache does at mmq-hopper-q1.cu:149-153.

    g_wcache_bytes += need;
    auto res = g_wcache.emplace(key, c);
    return &res.first->second;

---

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

[The-Monk]

Before the careful review starts: pushed one more commit (142d5268) with two prefill findings from PC-sampling mul_mat_q on gfx1201 — both HIP-path-only, both bit-identical by construction (prefill PPL reproduces exactly: Q1_0 11.6466, Q2_0 10.1465, wikitext ub-default ch4), so they're pure performance:

1. unpack_q1_0_bytes's HIP fallback was the expensive path. The fast __byte_perm unpack is #if !defined(GGML_USE_HIP), so AMD gets the per-bit (bit ? 1 : -1) ternaries — which compile to v_cmp_eq + v_cndmask chains serializing on VCC (PC-sampled at 9.5% of mul_mat_q), and whose live ranges push the mmq_x=128 tile to the 256-VGPR ISA cap with scratch spills in the hot loop (76 B/lane; 188 B on the need_check variant). Replaced with a borrow-free SWAR map: bit-spread to {0,1} bytes, then ((spread << 1) + 0x7F7F7F7F) ^ 0x80808080 maps 1→0x01 / 0→0xFF per byte with no carries and no VCC. On an equivalent tree this also eliminated the spills outright (vgpr 256→239, scratch→0) — and the obvious alternative, capping mmq_x at the spill-free 96 tile, measures 8–9% slower, so the SWAR keeps the large tile and drops the spills.

2. The q2 __byte_perm crumb unpack is correct on HIP but pays an emulation tax. HIP's __byte_perm has to convert PRMT's nibble-selector convention into v_perm_b32's control word at runtime, so each of the four perms costs several VALU ops here rather than one instruction. The SWAR equivalent (2-bit spread, then c−1 per byte carry-free) produces the same linear crumbs-minus-one ordering the perm/unshuffle pair computes; guarded HIP/MUSA, CUDA keeps the perms.

Measured (Bonsai-27B, 1× Radeon AI PRO R9700 gfx1201, ROCm 7.14, same-session A/B on this branch):

	before	after	Δ	our fork (same GPU)	PPL gate
Q1_0 pp2048	1252.5 ± 3.4	1327.6 ± 1.6	+6.0%	1332–1337	11.6466 exact
Q2_0 pp2048	1036.2 ± 26.2	1310.0 ± 1.8	+26.4%	1304	10.1465 exact
Q1_0 tg128	62.3	62.4 (unchanged)	—	68.0	—
Q2_0 tg128	46.8	46.8 (unchanged)	—	52.3	—

The "our fork" column is the roc10 tree these techniques were developed on, same R9700/ROCm 7.14 — with this commit the branch's prefill is now at parity with it. The remaining decode gap (62 vs 68 Q1_0, 47 vs 52 Q2_0) comes from launch-geometry and kernel-fusion tuning on that tree; those constants are kernel-implementation-specific and don't transfer by copy (we measured nwarps=8 cratering this branch's kernel to 8 t/s in the original PR round), so porting them properly is a future round rather than part of this one.

One question for anyone with NVIDIA hardware handy: the pre-existing (unguarded) q2 perm unpack feeds raw crumb pairs as PRMT selector nibbles. The ternary pair (2,2) produces selector value 10, which on NVIDIA engages PRMT's sign-replicate mode instead of a table lookup — I can't test CUDA here, but it looks like qx/qy could come out wrong for that (very common) crumb pair on the CUDA path. The HIP path no longer depends on it either way, but it's probably worth a 5-minute check with a Q2_0 perplexity run on an NVIDIA card before this merges.

Full run provenance (stack/firmware/clock telemetry) captured as bench cards; happy to attach if useful. The same two techniques are on their way upstream in generic form for mainline's ternary types (ggml-org#27127).

[The-Monk]

Follow-up on the upstream-to-mainline suggestion — we took the SWAR findings to mainline master and measured, and the case there turns out to be much bigger than on this fork.

Mainline baseline on AMD is running at less than half of what the types can do. On gfx1201 (Radeon AI PRO R9700, ROCm 7.14), ggml-org/llama.cpp master with the 27B Q1_0 model:

mainline master	pp2048	tg128
baseline	1024.6	29.46
+ SWAR HIP branches	1332.9 (+30.1%)	64.56 (+119%)

Root cause is the same __byte_perm story as here, but mainline has it in all four places (both vec_dots and both MMQ tile loaders) with no HIP consideration at all: HIP's __byte_perm emulates PRMT's nibble-selector convention with a runtime control-word conversion and has no mode bits, so the 10-perm Q1_0 unpack roughly halves AMD decode throughput. We probed the semantics empirically — HIP matches PRMT for selectors < 8 and diverges (mode-less wrap vs sign-replicate) at ≥ 8; by the structure of the duplicated LUT constants both endpoints happen to produce correct values, so mainline is correct-but-taxed on AMD, exactly like this branch was.

The fix is the same ~60-line, four-site SWAR patch as 142d5268 here, adapted to mainline's structure (its Q2_0 is group-64, and the loaders live in mmq-load-tiles.cuh). Ready-made branch, gated three ways on hardware:
The-Monk/llama.cpp branch hip-q1q2-swar, commit 61dc60335c — test-backend-ops MUL_MAT all supported q1_0/q2_0 cases pass; prefill perplexity bit-exact (11.6466); decode-path perplexity bit-exact (11.6361). CUDA path untouched. (We also A/B'd a fold-identity vec_dot variant — 2·dot(c,u) − sum(u) via the q8_1 sums — and it measured identical to the SWAR materialization on RDNA4, so the branch keeps the simpler bit-exact form.)

Our recommendation: this should go upstream under PrismML's name. You authored and upstreamed the types; the kernels are your architecture, and a doubling of AMD decode for them lands best coming from the type owners. Take the branch or the diff as-is, restructure freely, no attribution needed beyond whatever's natural. One heads-up for whoever submits: mainline's CONTRIBUTING.md AI policy requires the submitting author's own comprehensive manual review, an explicit disclosure of how AI was used (these kernels were developed and A/B-tested with AI assistance on our hardware), and a personally-written PR description. If you'd rather we submit it ourselves with you reviewing, that works too — but our preference is that it ships under the flag of the people who built the types.

[The-Monk]

Completing the mainline table with Q2_0 — we requantized the 27B ternary model to mainline's group-64 Q2_0 (via F16; near-lossless for ternary values, PPL 10.1383 vs 10.1465 on this fork's g128) and ran the same A/B on master vs the hip-q1q2-swar branch:

mainline, gfx1201	pp2048	tg128	PPL (wikitext ch4)
Q1_0 master	1024.6	29.46	11.6466
Q1_0 + SWAR	1332.9 (+30.1%)	64.56 (+119%)	11.6466 bit-exact
Q2_0 (g64) master	1066.2	33.47	10.1383
Q2_0 (g64) + SWAR	1320.7 (+23.9%)	43.92 (+31.2%)	10.1383 bit-exact

So the branch is fully characterized on both types: four bit-exact gates, test-backend-ops green, CUDA path untouched. On RDNA3: the guard is GGML_USE_HIP-wide and the SWAR is pure integer math with no cross-lane ops, so correctness is architecture-independent; the __byte_perm emulation cost lives in ROCm's headers identically on every AMD arch, so we'd expect the gains to transfer — your RDNA3 validation would turn "measured on RDNA4" into "validated across consumer RDNA", which is exactly the evidence a mainline reviewer will want. Happy to share the g64 GGUF if a matching test model saves you time.

[The-Monk]

One more dataset for the upstream decision, since we hold the AMD hardware in this collaboration: we measured the full ternary scale-granularity curve on Bonsai-27B — including mainline's TQ2_0, which is byte-for-byte the "g256" point — so the format question for mainline distribution can be decided on data rather than taste. Everything below is from our instrumented rig (1× Radeon AI PRO R9700 gfx1201, ROCm 7.14; methodology notes at the end).

Quality vs scale granularity (same model content, requantized per format)

format	group	bpw	PPL ch20 ± SE	vs g128	verdict
mainline TQ2_0	256	2.0625	11.4606 ± 0.458	−0.9%, best measured	quality-viable, not merely compat — but zero GPU kernels upstream
mainline Q2_0	64	2.25	11.5574 ± 0.472	≈ tie	+9% bytes over TQ2_0 for nothing
fork Q2_0 (yours)	128	2.125	11.5616 ± 0.473	baseline	the fork optimum; matches your historical 11.5603
experimental g1024	1024	2.0156	12.2978 ± 0.521	+6.4%	dominated — see below

Exception-tensor accounting, since it matters: the g64 and TQ2_0 files carry q4_K/q6_K token-embd/output (mainline's rule), the g128 and g1024 files carry the fork's q2_0 ones. The decisive cell is therefore TQ2_0 vs g64 at identical exceptions: −0.84% for TQ2_0 — the advantage is the format/quantizer itself. (ch4 quick-pass numbers told the same ordering but sit inside their ±0.95 SE; ch20 is the number to trust. CPU-vs-GPU evaluation cross-checked on g128: within ±0.01.)

Why very coarse groups lose — and why g256 doesn't. We analyzed the shipped g128 GGUF's scale structure directly (per-tensor, per-row block statistics over all 498 ternary tensors):

• per-tensor scale rel-std: 13.5% mean, 29.9% max
• split across rows: 10.9% mean / 29.0% max — real per-channel magnitude structure
• split within rows: 7.2% mean / 14.1% max — also real, and this is what coarse groups flatten

Amax-based coarse scales zero out survivors in low-magnitude groups (any group whose local scale sits under ~half the merged amax loses its ±1s to rounding). Merging two g128 groups (→g256) stays inside the tolerable band; merging eight (→g1024) crosses it and costs +6.4% even with an MSE-optimal magnitude second pass (d = mean |w| over survivors, codes preserved) — the within-row variance is trained signal. We built the g1024 type end-to-end (CPU+HIP kernels, two-pass quantizer) so that number is a measurement, not an extrapolation; it also ran 1.3% slower than g128 at bs=1 because the byte savings don't survive untuned launch geometry. The quality plateau runs from g64 through g256 and falls off a cliff before g1024 — and TQ2_0 sits at the cheap end of the plateau.

AMD performance by format (same GPU, post-SWAR kernels everywhere)

format	tree	pp2048	tg128	notes
Q2_0 g128	this fork @ 142d5268	1310.0	46.8–52.4*	*decode varies with the session's thermal state; interleaved A/Bs used throughout
Q2_0 g128	our perf fork	1332	52.3	adds swept launch geometry (rpb=3, nwarps=8)
Q2_0 g64	mainline + SWAR branch	1320.7	43.92	−6–7% decode vs g128, matching the +5.9% byte cost on a memory-bound path
Q2_0 g64	mainline master	1066.2	33.47	the perm-emulation tax, for scale
TQ2_0 g256	mainline	—	—	zero GPU kernels upstream (0 references in ggml-cuda); CPU-only today
TQ2_0 (projected)	with the scoped port	~1330	~53	byte model over tuned g128-class kernels (2.0625/2.125 × 52.3); the fold + SWAR kernels map 1:1 onto its layout

Roofline context for the decode column: our measured pure-read ceiling on this card is 640.0 GB/s (custom streaming-kernel calibration, reproduced ×2 — we use measured, not datasheet), and the g128 decode kernel runs at ~87% of it with exact per-model byte accounting from GGUF metadata — i.e., ternary decode on RDNA4 is memory-bound, which is why bytes-per-weight dominates every row of that table and why sub-percent bpw differences show up in tg.

What this means for the upstream/distribution question

1. TQ2_0 is the data-supported mainline distribution format: equal-or-better quality than both Q2_0 layouts, the fewest bytes (−3% vs g128, −9% vs g64), already upstream-native — no layout-convergence fight needed. g128 remains the right fork format (its GPU kernels are mature and its decode is fastest today).
2. The blocker is that TQ2_0 has zero GPU kernels upstream — CPU-only. Our fold-identity + SWAR kernels map 1:1 onto its layout (same c−1 ternary math, per-256 scale), and we're happy to supply that port as part of the upstream package; projected AMD decode from the byte model is ~53 t/s on this card (vs 52.3 for g128 on tuned kernels), i.e. TQ2_0 would also be the fastest once ported.
3. The full provenance for every number here (stack versions, firmware, clock/junction telemetry per run, bit-exact PPL gate values, kernel-level counter data) is captured as bench cards from our provenance framework — we can attach the JSON cards or give you the raw sweep logs if useful for the PR.

Methodology: PPLs are wikitext-2, 4×512 chunks, prefill-path unless marked; quality comparisons use requants of the identical g128 source through F16 (lossless for ternary values); TQ2_0's PPL is CPU-evaluated (no GPU path exists) with a same-config CPU-vs-GPU calibration run on g128 agreeing within ±0.01 PPL. Decode A/Bs are same-session interleaved with VRAM-drain guards; perf numbers are r=3–5 llama-bench with reported ±.

[The-Monk]

Follow-up from the AMD measurement series — this one is a capability proposal rather than a kernel finding.

Proposal: ship a nextn/MTP head for Bonsai-27B. We validated MTP self-speculation on AMD (llama.cpp --spec-type draft-mtp): an fp8 Qwen3.6-27B with its stock head gains ×1.63 effective decode at 43% accept. We then grafted that same head onto Bonsai-27B-Q1_0 (vocab is identical — 248,320 token-for-token — and the graft loads and runs cleanly at decode parity) and measured 32.7% accept: a real ancestry signal, but below breakeven on a ternary target whose base decode is already ~67 t/s on RDNA4.

A Bonsai-distilled head (1 transformer layer + eh_proj/norms, ~400 MB bf16/fp8, trained against Bonsai hidden states) should recover the ~50%+ accept that makes speculation pay on fast ternary targets; projected ~1.5× effective decode on top of all the kernel wins, on every GPU vendor — spec decode multiplies whatever the platform baseline is.

We can supply the gguf graft tooling (drop-in file format already worked out) and same-day AMD benchmarks for any candidate head: llama-speculative-simple --spec-type draft-mtp, accept + effective-t/s sweep at draft depths 2/3/4.

[The-Monk]

Update — measured result worth flagging: the stock Qwen3.6 MTP head is already net-positive on Bonsai-27B Q2_0, no training needed.

target	accept (pos-1)	effective t/s	vs plain
Bonsai-27B Q2_0 (graft, n-max=2)	49.2%	62.9	+20.9% (52.0 base)
Bonsai-27B Q1_0 (graft, n-max=2)	32.7%	61.4	−8.5% (67.1 base)

(RDNA4 R9700, temp 0 — greedy verify is exact-match, so output is bit-identical to plain greedy decode.)

The split is informative: Q2_0 preserves the ternary master faithfully and the head recognizes it; Q1_0s binarization drifts the distribution and accept drops below breakeven on a 67 t/s target. Which sharpens the earlier proposal: a Bonsai-distilled head should lift Q1_0 into the same ~50% accept band (projected ~67 → 85+), and likely push Q2_0 further still. Meanwhile the Q2_0 +21% is available to your users today — the graft is 15 tensors (blk.64.* from any Qwen3.6-27B MTP gguf) + block_count 65 + nextn_predict_layers=1; happy to share the script.

[The-Monk]

Pushed one more commit to the branch: AVX2 tier for the Q2_0 CPU vec_dot.

The current fast path is gated on VNNI (__AVXVNNI__ / AVX-512-VNNI), so every pre-2021 x86 — Haswell through Zen 3 — silently runs the scalar loop. The new #elif __AVX2__ tier uses a pshufb-replicate + nibble-LUT crumb decode and maddubs with the same dot(c-1,u) = dot(c,u) − sum(u) fold as the VNNI tier.

Measured (Bonsai-8B Q2_0, dual Xeon E5-2699 v4 = AVX2-only, 11 threads): scalar 1.81 t/s → 5.90 t/s (3.26×), PPL bit-exact (10.4508 both ways, 2-chunk gate).

One observation on the existing VNNI tier while we were in there (not changed — we have no VNNI hardware to validate on): it horizontal-sums both dpbusd results to scalar every 32 elements (hsum_i32_8 ×2 per block). Keeping a vector accumulator and folding d0*d1 via FMA — the structure the new AVX2 tier uses — should save ~15–20 ops per 32 weights there too, if someone with VNNI silicon wants to verify.

Thread-scaling note for CPU users: on the 2-socket test box, t=11 beat t=44 (6.0 vs 4.5 t/s) — per-op sync dominates at high thread counts on many-core machines; worth a mention in docs.

[The-Monk]

CPU follow-up to the AVX2 vec_dot tier from this morning — two more commits on the branch, closing out the x86 story for older cores.

The gap: the 4x8 repack path for Q1_0/Q2_0 (ggml_gemv/gemm_*_4x8_q8_0) was AVX-512-VNNI-gated on x86, so everything Haswell through Zen 3 fell back to unbatched vec_dot — and on dual-socket parts the per-op sync cost made decode inverse-scale past ~22 threads.

c6e2a0c — AVX2 tier for the 4x8 repack kernels (+ AVX2 engagement). Weights expand to unsigned bytes ({0,1} bits / {0..3} codes) so the u8×s8 dot runs as maddubs+madd — pair sums bound at 768, no s16 saturation; col-pair-packed fp32 accumulators keep the 4-row GEMM tile inside AVX2 16-ymm budget.

4ad267e — opt-in GGML_NUMA_MIRROR=1. Per-NUMA-node copies of repacked weight tensors (mmap+mbind, no libnuma dependency); each thread reads the copy local to the cpu it is currently on (sched_getcpu per chunk), so locality survives thread migration and the chunk scheduler needs no changes. Copies are made post-repack → bit-exact.

Measured on dual E5-2699 v4 (2016 Broadwell — AVX2, no VNNI), Bonsai-8B Q2_0, same-session A/Bs:

tg128 t/s	t=11	t=22	t=44
before (vec_dot path)	6.5	7.6	5.8
repack AVX2	10.6	17.3	19.8
+ numactl --interleave=all --no-mmap			26.2
+ GGML_NUMA_MIRROR=1		20.4	27.0

pp512: 10.6/21.1/39.1 → 23.0/45.7/64.9. Bonsai-27B Q1_0 tg64: 2.6 / 2.5 / (collapse) → 3.6 / 5.8 / 6.5. Known tradeoff: Q1_0 prefill is ~13% slower than its (already strong) vec_dot path — the q1 gemm bit-expand is untuned; decode dominates the CPU use case so engagement is kept for both types.

Gates: PPL parity for both types (Δ0.003 ≪ SE), mirror bit-exact, 54-shape kernel harness vs the generic reference (max rel err 4e-6, fp-order only).

One diagnostic worth recording: with mirroring, weights are fully node-local and decode moved only +3% at t=44 (~50 GB/s of ~140 available) — post-repack, ternary CPU decode is not bandwidth-bound. Profile puts the Q2 gemv at IPC≈1 with the expand shuffle-port pressure as the likely wall; that (not data placement) is the next lever if anyone wants to chase it.

Net: 5.8 → 27.0 t/s decode (4.7×) on 2016 hardware, one day of AVX2 work. HT hurts decode (t=88 craters); pin to physical cores.

[The-Monk]

Follow-up with Bonsai-27B numbers, same dual E5-2699 v4 (2016 Broadwell — AVX2, no VNNI), after four more commits on the branch: 4ad267e5 (opt-in GGML_NUMA_MIRROR=1), 8b6271d3 (lock-free mirror lookup), 842b1b1c (gemv scale hoist), d90dab7f (odd-vocab repack tails).

tg64 t/s (Q1_0 / Q2_0)	t=11	t=22	t=44
before (vec_dot path)		3.33 / 2.59	2.30 / 2.27 (inverse scaling)
repack AVX2	3.88 / 3.12	6.26 / 5.01	7.28 / 5.20
+ GGML_NUMA_MIRROR=1	4.01 / 3.22	6.43 / 5.35	8.23 / 6.96

Net at 44 threads: Q1_0 2.30 → 8.23 (3.6×), Q2_0 2.27 → 6.96 (3.1×), and the dual-socket inverse-scaling is gone for both formats. Mirroring is worth +13% / +34% (Q1/Q2) here versus the ~+3% it showed on the 8B — the bigger the weight set, the more node placement matters.

Two notes for anyone benchmarking multi-socket CPU decode:

1. The vec_dot baseline is placement-lottery-sensitive: the same command measured 2.46 and 3.33 t/s (Q1_0, t=22) hours apart, tracking how the mmap page cache happened to land across NUMA nodes after unrelated I/O. The mirror removes that variance entirely — decode speed becomes placement-independent.
2. d90dab7f matters beyond this model: any vocab not divisible by 4 (e.g. qwen3's 151669) previously excluded output.weight — the largest single decode matmul — from the repack path. Aligned rows tile-repack as before; the ≤3 tail rows stay in original block format at the end of the buffer and route through the type-generic vec_dot. PPL bit-identical to the vec_dot reference.

[The-Monk]

Round-4 MTP heads are trained and scored — the KL + chain-unroll recipe works, and Q1_0 self-speculation now reaches 95 t/s on code.

Recipe delta vs the round-2/3 heads posted earlier: 30MB domain-weighted self-generated corpus per format (code/math/technical-heavy, llama-batched -np 32), capture upgraded to record top-8 target logits per position (KL ranking-distillation term, weight 0.5), and a two-step chain unroll in the trainer (the head consumes its own post-head-norm hidden state as the next-step input — matching drafting mechanics, weight 0.5). 4000 steps from the prior-round init, frozen base embd/output as before.

Bonsai-27B, 2× R9700 (numbers below from one GPU, same-session; 512-token greedy, temp 0):

t/s (accept)	technical	code	creative
Q1_0 plain decode	66.1	66.1	66.1
Q1_0 + r4 head, n=2	78.8 (56.0%)	91.8 (77.2%)	75.0 (50.8%)
Q1_0 + r4 head, n=3	75.0 (45.8%)	95.2 (70.9%)	71.3 (41.8%)
Q2_0 plain decode	51.7	51.7	51.7
Q2_0 + r4 head, n=2	64.5 (51.0%)	73.9 (67.7%)	58.8 (41.0%)
Q2_0 + r4 head, n=3	61.2 (40.1%)	77.5 (62.8%)	52.5 (29.8%)

Findings worth recording:

1. Depth 3 now beats depth 2 on code for both formats — first time in this series. The chain-unroll term did exactly what it was added for: Q1 chained accept at n=3 went 50.7% → 70.9% vs the round-2 head. Depth 4+ still loses (85.9 / 82.6 t/s at n=4/6) — chained accept decays too fast past 3.
2. The domain weighting overshot: code improved massively, but technical accept regressed ~4-6 pts vs the prior heads on both formats. Per-workload head choice is currently the honest deployment answer; a rebalanced corpus is the next round.
3. 2:4-sparse head experiment (three training rounds): pruning the head to 2:4 fp8 (block_2of4_fp8) costs a structural −5-6 pts accept that survived 4× recovery training — but the head-only artifact drops ~470MB → ~310MB. Worth it for downloads, not for peak speed.
4. Q1_0 + r4 head at n=2/3 is faster than Q2_0 plain by ~40-80% at Q1 file size (3.5GB vs 6.7GB) — but Q2_0+head still holds the quality/speed crown from the earlier finding (better PPL at comparable effective speed).

Head-only grafts (~470MB, apply with the graft tool already on the release branch) available on request — same drop-in format proposed earlier in this thread.

[bri-prism]

Thanks for all of this, and for the standing offer on the R9700 box. The RDNA4 data is exactly what we were missing: our AMD work so far has been on MI300X only.

To get the kernel part merged quickly, a few requests:

1. Rebase onto prism-v7. prism is retired (see the README); prism-v7 is the active branch and is where releases come from. The PR currently shows +14.8k lines and a conflict because the branch predates a lot of prism, so a rebase will also shrink the diff to the 11 files you actually changed. Two things to know about v7 when you rebase:

• Q2_0 on v7 is mainline's group-64 layout (id 42 is now the upstream block), and our group-128 ternary block is PQ2_0. Your Q2_0 vec_dot will need to be applied to both.
• hip: faster Q1_0, Q2_0 and PQ2_0 decode on AMD #125 landed yesterday with a HIP path for the same three dot products (Q1_0, Q2_0, PQ2_0): it uses the hardware byte permute (__builtin_amdgcn_perm, one v_perm_b32 per 4 codes) for the unpack and keeps the split-K MMVQ geometry on CDNA. On MI300X it is 2.9 to 3.7x on Q1_0 decode and 1.9 to 2.6x on PQ2_0, KL against the base build is 0. Your SWAR + dot(s,u) = dot(c,u) - sum(u) formulation is a different route to the same place, so please measure the two against each other on gfx1201 after the rebase rather than replacing one with the other. If yours wins on RDNA4, an #if on the arch is fine; if they tie, fewer code paths wins. We can run the same A/B on MI300X, the pod recipe is ready.

2. Split the PR into three. They have very different review costs and it would be a shame for the vec_dot change to wait on hipBLASLt:

• (a) HIP vec_dot for Q1_0 / Q2_0 / PQ2_0 plus the unpack_q1_0_bytes SWAR for the MMQ prefill path. Small, default-on, bit-exact. This can land first.
• (b) Opt-in activation-quant dedup. Small, but see the stream-safety point below.
• (c) Opt-in hipBLASLt prefill routes and the weight cache. 1,300 new lines and a new optional dependency; this one needs the most review and will take longest.

3. Address the Copilot findings before re-requesting review. Most of them are real:

• vecdotq.cuh signed shift UB: you fixed the Q2_0 pack; please check the Q1_0 path and MMQ unpack for the same pattern.
• Dedup cache hits are not stream-safe. Sibling matmuls in concurrent graph regions run on different streams, and a hit consumes the buffer without waiting on the quantization that produced it. Record an event on the miss stream and wait on it at hit, or restrict sharing to the same stream.
• hipBLASLt weight cache keyed by raw device address: needs at least (device, address, N, K) and the device passed through invalidation. The Hopper cache in mmq-hopper-q1.cu shows the pattern.
• Support predicates must check strides (non-contiguous activation views currently produce wrong results instead of falling back), and the Q1_0 route needs the same free-VRAM preflight the Q2_0 route has.
• The E4M3 rounding at the subnormal boundary.

4. PR description. The description carries a "Generated with Claude Code" footer. AGENTS.md in this repo (lines 45 to 47, inherited from upstream) lists AI-written PR descriptions under prohibited usage, so please rewrite it in your own words when you rebase; it can be much shorter now that the numbers are in the thread.

The other threads here (the scale-granularity curve, the MTP head grafts and training recipe, the AVX2 CPU tiers, NUMA mirroring) are each valuable and each deserves its own issue or PR so they can be reviewed on their own terms; folding them into this branch is what pushed it to 58 commits. The CPU AVX2 work in particular looks like a clean standalone PR.

On upstreaming: agreed with @khosravipasha, the +119% Q1_0 decode on mainline master is a big deal and it is your finding, so please open that PR on ggml-org yourself. #125 is CDNA-scoped, and once the two approaches have been compared on both architectures we can point the upstream PR at whichever is faster on each.

[bri-prism]

MI300X numbers for the A/B I offered above, so the arch question is settled before you rebase.

Setup: prism-v7 at 3c12dbd, three builds of one tree that differ only in the HIP dot product for Q1_0 / Q2_0 / PQ2_0. A is #125 as merged (one v_perm_b32 LUT per 4 codes). B is your formulation ported onto v7: bit-spread the raw codes into byte lanes, dp4a on {0,1} or {0..3} directly, and apply the offset once per 32-element chunk through the q8_1 stored sum (2*dot(c,u) - sum(u) for Q1_0, dot(c,u) - sum(u) for the 2-bit types). C is A with the nibble spread done by fused shift-or instead of a 32-bit multiply, to isolate that difference. ROCm 6.4.1, gfx942, llama-bench -ngl 99 -fa 1 -n 128 -r 5, rounds A B C then C B A.

model	A: v_perm LUT (r1 / r2)	B: SWAR + identity (r1 / r2)	C: v_perm, shift-or (r1 / r2)	B vs A
4B Q1_0	317.0 / 317.1	324.4 / 324.9	315.7 / 317.0	+2.4%
8B Q1_0	265.9 / 265.7	283.1 / 281.9	266.6 / 265.5	+6.3%
8B PQ2_0	253.5 / 255.6	259.9 / 258.9	255.8 / 254.9	+1.9%
27B Q1_0	97.4 / 97.4	102.4 / 103.1	97.8 / 97.7	+5.4%
27B PQ2_0	91.6 / 91.7	93.5 / 93.9	91.8 / 91.8	+2.2%

Your formulation wins on CDNA as well, on every model, and the round-to-round spread is under 0.7 tok/s so none of it is noise. Q1_0 gains more than the 2-bit types because it was issuing four permutes per 16 weights against two. C equals A, so the multiply versus shift-or spread is not where the gain is; skipping the symbol materialization is.

Correctness: all three builds pass test-backend-ops MUL_MAT (135 OK, 0 FAIL) and MUL_MAT_ID (225 OK, 0 FAIL) for the three types. B is not bit-identical to A by construction (the offset is applied in float once per chunk instead of inside the integer dot), but scoring B against A's logits over 60 chunks of wikitext-2 gives mean KLD 0.000000, max 6.4e-5 and 99.99% same top token, which is the resolution floor of the tool. Prefill is unchanged (4956 vs 4968 pp2048 on 8B Q1_0), as expected since the MMQ tiles are not touched.

So: one AMD path, yours, for both CDNA and RDNA. When you rebase the vec_dot piece onto prism-v7, please replace the v_perm helpers from #125 rather than adding an arch switch; the port I used is 60 lines and I am happy to share it if it saves you time. The same trick for the MMQ load_tiles unpack on the prefill side is untested on MI300X; if you include it in the same PR I will run it there too.

[The-Monk]

Thanks for running that A/B, and for including the C arm — that's the control that actually settles it, since it rules out the multiply-vs-shift-or difference and leaves skipping the symbol materialization as the cause. Your numbers plus mine point the same way on both architectures, so I'll take you up on the port. Drop it wherever is easiest (branch here, patch, or gist) and I'll rebase on top of it rather than redo the 60 lines.

Plan for the first PR, matching your split:

• HIP vec_dot for Q1_0 / Q2_0 / PQ2_0, replacing the v_perm helpers from hip: faster Q1_0, Q2_0 and PQ2_0 decode on AMD #125 outright rather than adding an arch switch.
• Q2_0 applied to both the group-64 block now at id 42 and to PQ2_0.
• unpack_q1_0_bytes SWAR on the MMQ prefill path included in the same PR, so you can run that one on MI300X — it's untested there and I only have RDNA4 here.
• Signed-shift UB swept across the Q1_0 path and the MMQ unpack, not just the Q2_0 pack I already fixed.
• Description rewritten in my own words.

I'll re-measure A vs B on gfx1201 after the rebase and post those numbers next to yours before I re-request review, so the RDNA4 side is on the record the same way.

Dedup and hipBLASLt follow as separate PRs: dedup with the completion-event fix (event on the miss stream, wait at hit), hipBLASLt with the cache keyed on (device, address, N, K) and the device passed through invalidation, stride checks in both support predicates, the free-VRAM preflight on the Q1_0 route to match Q2_0, and the E4M3 subnormal-boundary rounding.

Agreed on upstream — I'll open the ggml-org PR myself, but after the vec_dot lands here so the two don't diverge while it's in review.

The R9700 offer stands, and the standing invitation to run anything of yours on it does too.

[The-Monk]

Ran the A/B on gfx1201, and I have to correct something before anything else: I nearly sent you a message saying my decode formulation does not win on RDNA4. That was wrong, and it was wrong because I ran a two-arm test and drew a three-arm conclusion.

Measuring mine against #125 shows a wash. That is equally consistent with "mine stopped working" and with "#125 independently captured the same win," and I had no arm that separated them. So I built one: prism-v7 at 3c12dbd with the three vec_dot HIP guards forced to #if 0, falling through to the original __byte_perm path. Same commit, same tree, same ROCm for all three arms, ldd-verified identical runtime libs.

Decode (tg128, 27B, -fa 1 -ngl 99 -n 128 -r 5, two rounds, arm order rotated)

model	BYTEPERM (baseline)	PERM (#125)	SWAR (mine)	#125 vs base	mine vs base	mine vs #125
27B Q1_0	29.27	64.94	64.22	+121.8%	+119.4%	-1.1%
27B Q2_0 (g64)	33.33	42.92	43.71	+28.8%	+31.1%	+1.8%
27B PQ2_0	33.69	48.85	48.52	+45.0%	+44.0%	-0.7%

Round-to-round spread 0.33-1.06 t/s. Both routes fix the same defect and both are worth 1.3x to 2.2x over the __byte_perm path. Your +119% figure for Q1_0 on mainline reproduces here at +119.4% on a different branch with different models, which is a good independent check on both our measurements.

On RDNA4 the two routes are indistinguishable: -1.1%, +1.8%, -0.7%, each inside its pair's round-to-round spread. That is a wash, not a regression -- and combined with your CDNA numbers it points where you already said it does. Yours measured +2.4% to +6.3% for the SWAR route on gfx942; mine measures no cost for it on gfx1201. So unifying on SWAR is free on RDNA4 and worth 2-6% on CDNA, while unifying on v_perm is free on RDNA4 and gives up that 2-6%. Same number of code paths either way, so "fewer paths" does not decide it -- the CDNA margin does.

Conclusion: unify on SWAR, no arch #if -- which is exactly what you proposed. I will replace the v_perm helpers from #125 rather than adding a switch, and I will take the port you offered.

One thing worth flagging whichever way this lands: #125 is not arch-gated. Every guard in vecdotq.cuh is #if defined(GGML_USE_HIP) with no arch macro, so it currently ships on RDNA2, RDNA3, RDNA3.5, RDNA4 and CDNA, and has been benchmarked on two of them (your MI300X, my gfx1201). Whatever we unify on inherits that reach, so an #if drawn on architecture would silently assign RDNA3/3.5/RDNA2 to a branch nobody has measured. That is an argument for one path, not two.

Prefill (pp2048) — this is the part that is actually mine

model	BYTEPERM	PERM (#125)	SWAR+MMQ (mine)	#125 vs base	mine vs base
27B Q1_0	1033.4	1026.6	1307.0	-0.7%	+26.5%
27B Q2_0 (g64)	1053.6	1052.7	1292.4	-0.1%	+22.7%
27B PQ2_0	1052.5	1052.8	1295.4	+0.0%	+23.1%

#125 is flat on prefill, as expected since it only touches the decode vec_dot. The MMQ load_tiles unpack is worth +23-27% pp2048 on all three types and nothing on prism-v7 currently has it.

Mechanism, from the compiled kernel descriptor (mmq_x=128, llvm-readobj on the unbundled gfx1201 code object):

kernel	before	after
mul_mat_q<Q1_0,128>	256 VGPR, 80 B/lane scratch	212 VGPR, 0 B
mul_mat_q<Q2_0,128>	256 VGPR, 88 B/lane scratch	227 VGPR, 0 B
mul_mat_q<PQ2_0,128>	256 VGPR, 88 B/lane scratch	227 VGPR, 0 B

The tiles were spilling; removing the spills is the win. Unlike the decode path this one is bit-identical, not merely close: I compared the SWAR unpack against ROCm's __byte_perm across the full 16-bit input space, 0/65536 mismatches for both the Q1_0 4-value and the 2-bit 4-value unpack. test-backend-ops MUL_MAT and MUL_MAT_ID pass with 0 failures.

Both of the obvious objections are closed by measurement. An MMQ-only arm (decode vec_dot left as #125's) reproduces the full win on all three models to within 0.05% -- the decode change contributes nothing to prefill, which static analysis confirms structurally (the vec_dots are dispatched only from mmvq.cu, never from the MMQ path). And the win is flat across pp512/1024/2048/4096 (+26.3/+27.3/+27.2/+26.3% on Q1_0; +22-23% on both 2-bit types at 2048/4096), so it is not a single tile-regime artifact.

One thing your A/B and mine cannot resolve between them

You measured on gfx942 with ROCm 6.4.1; I measured on gfx1201 with 7.14. Two variables differ, so neither of us can attribute the small decode differences to architecture. It does not change the decision -- SWAR is at worst free on RDNA4 and better on CDNA either way -- but if you ever want that question settled, re-running your A/B once on 7.x would do it -- the SWAR exists to route around __byte_perm lowering badly on HIP, and that is a compiler property that may simply have been fixed.

RDNA4 launch geometry: an additional +8% decode on the 2-bit types, through the existing arch table

Separate from the unpack question entirely: calc_nwarps/calc_rows_per_block have no RDNA4 entries for these types, so they launch on defaults. Sweeping them on gfx1201 gives PQ2_0 rpb=2/nwarps=6 and Q2_0(g64) rpb=2/nwarps=6; Q1_0's stock config is already its peak (decline confirmed monotonically through nwarps 1..8).

Measured with the tuning applied on top of the SWAR unpack (tg128, 27B, two rotated rounds):

type	untuned	tuned	gain
PQ2_0	48.27	52.30	+8.4%
Q2_0 (g64)	44.24	47.85	+8.2%
Q1_0	64.72	64.49	wash

Also verified on top of the v_perm unpack (52.32 / 48.09) -- the constants are robust to the unpack choice, so this composes with whichever dot product you keep. It goes through MMVQ_PARAMETERS_RDNA4, the arch-keyed mechanism mmvq.cu already uses, so the arithmetic stays one path.

It rides with a correctness fix you want regardless. Enabling rpb>1 exposed that the mul_mat_vec_q row-bounds guard compares row0 + i against stride_col_dst, which under MUL_MAT_ID is the expert-channel stride, not the row count. rpb>1 is already reachable today through the GENERIC/GCN/TURING/GB10 tables at ncols_dst 2-8 (ordinary batched MoE), so any expert tensor whose row count is not divisible by rpb writes out of bounds on those backends now, on your tree as it stands. Reproduced with m=517, n=1 (PQ2_0/Q2_0/Q4_K fail, then pass with the fix). Worth noting the stock test-backend-ops MUL_MAT_ID matrix never catches this -- its m values (64, 512) divide evenly by the rpb the tables pick, so I added the m=517 case alongside the fix. I would put the guard fix up as its own small PR since it is live for non-RDNA4 backends independent of any tuning.

The split, as you asked for it

• (a) HIP vec_dot for Q1_0 / Q2_0 / PQ2_0 replacing the hip: faster Q1_0, Q2_0 and PQ2_0 decode on AMD #125 v_perm helpers, plus the MMQ load_tiles unpack, plus the signed-shift UB fix. Small, default-on, and the MMQ half is bit-exact. Your offer to run the MMQ half on MI300X is what I would most like to take you up on -- register pressure does not transfer between architectures, and the spills it removes on gfx1201 may or may not exist on gfx942.
• (b) activation-quant dedup, with the completion-event fix for the stream-safety issue.
• (c) hipBLASLt prefill routes, with the (device, address, N, K) cache key, the free-VRAM preflight on the Q1_0 route, and the E4M3 boundary fix.
• (d) the MUL_MAT_ID row-bounds guard fix (small, correctness, live on your tree today), and on top of it the RDNA4 mmvq table entries above.

On the Copilot signed-shift item: qs is uint8_t, so b << 21 tops out at 0x1FE00000 and the Q1_0 path and MMQ unpack are already clean. The genuine instance was the Q2_0 byte pack, which is fixed. The stride checks, cache key, VRAM preflight and E4M3 rounding are all real and will be in (b)/(c).