Skip to content

HIP: Tune RDNA4 MMQ tiles for Q6_K / Q2_K (gfx1201) - #25587

Closed
qiongz wants to merge 3 commits into
ggml-org:masterfrom
qiongz:hip-rdna4-mmq-mmvq-mtp-opt
Closed

HIP: Tune RDNA4 MMQ tiles for Q6_K / Q2_K (gfx1201) #25587
qiongz wants to merge 3 commits into
ggml-org:masterfrom
qiongz:hip-rdna4-mmq-mmvq-mtp-opt

Conversation

@qiongz

@qiongz qiongz commented Jul 12, 2026

Copy link
Copy Markdown

Overview

This PR aims to tune mul_mat_q on RDNA4 (gfx1201). Changes are gated with #if defined(RDNA4). Non-RDNA4 builds match upstream. Benchmarked on Radeon AI PRO R9700 / ROCm 7.2.4; shapes from Qwen3.6-27B FFN.

Main Results

mmq-bench prefill ( ffn_gate 17408x5120, same-GPU0 A/B)

type stock n512(TFLOPS) this-version n512(TFLOPS) delta stock n2048(TFLOPS) this-version n2048 (TFLOPS) delta
Q6_K 32.64 46.75 +43.3% 33.04 46.64 +41.2%
Q4_K 76.32 76.50 +0.2% 76.15 76.37 +0.3%
Q4_0 76.75 77.89 +1.5% 76.59 77.93 +1.7%
Q5_K 66.80 71.43 +6.9% 66.79 71.23 +6.6%
Q5_0 73.57 76.25 +3.6% 73.25 76.04 +3.8%
Q8_0 80.07 80.68 +0.8% 81.01 81.29 +0.3%
Q2_K 1.93 6.50 +237% 2.07 6.51 +214%
Q3_K 58.10 58.19 +0.2% 58.12 58.74 +1.1%

27B Q6_K prefill (Qwen3.6-27B-UD-Q6_K_XL)

test stock t/s this version t/s delta
pp512 705.83 888.79 +25.9%
pp2048 694.68 872.70 +25.6%
pp8192 641.65 792.90 +23.6%
tg128 21.50 21.50 +0.0%

27B Q2_K prefill (Qwen3.6-27B-UD-Q2_K_XL)

test stock t/s this version t/s delta
pp512 87.21 263.99 +202.7%
pp2048 87.10 262.42 +201.3%
pp8192 86.64 254.39 +193.6%
tg128 30.80 30.80 +0.0%

llama-bench 7B Q6_K (Qwen2.5-7B-Instruct-Q6_K)

test stock t/s this version t/s delta
pp512 2279.14 3289.76 +44.3%
pp2048 2302.51 3297.88 +43.2%
pp8192 2145.20 3000.91 +39.9%
tg128 85.19 85.59 +0.5%

Decode

E2E MTP accept(Qwen3.6-27B-UD-Q6_K_XL)
n-max metric stock (e3546c7) this version (fe2df7f) delta
N2 accept 86% (60/70) 86% (60/70) identical
N2 t/s 42.64 ± 0.09 42.40 ± 0.08 -0.6%
N3 accept 77% (66/86) 77% (66/86) identical
N3 t/s 44.88 ± 0.07 44.55 ± 0.07 -0.7%
Q2_K/Q3_K decode regression fix

Q2_K/Q3_K decode (N=14/32/128): no regression vs stock. Q2_K N=128 was -70% with intermediate build, fixed to -0.2% in final commit.

Perplexity Test

llama-perplexity 7B (wiki.test.raw): stock 7.9666, opt 7.9675 (+0.011%, limit 0.5%)

Test-backend-ops

test-backend-ops MUL_MAT: 1134/1134 passed (ROCm0 backend)

Additional information

Precision (vs stock golden, --check-ref)

type N=14 N=128 N=512 N=2048
Q2_K bit-exact bit-exact near 6.9e-5 near 6.9e-5
Q3_K bit-exact bit-exact bit-exact bit-exact
Q4_K bit-exact bit-exact bit-exact bit-exact
Q4_0 bit-exact bit-exact bit-exact bit-exact
Q5_K bit-exact bit-exact bit-exact bit-exact
Q5_0 bit-exact bit-exact bit-exact bit-exact
Q6_K bit-exact near 3.1e-5 near 3.8e-5 near 3.8e-5
Q8_0 bit-exact bit-exact bit-exact bit-exact

Summary: 27/32 BIT-EXACT, 5/32 NEAR (cos=1.0, argmax=1.0, abs<1e-4). No FAIL.
NEAR causes: Q6_K N>=128 = fmaf epilogue; Q2_K N>=512 = mmq_x=64 vs stock 128 reduction order.

Final mmq.cuh state (what's gated on type, RDNA4 only)

Lever Q6_K Q2_K Q3_K Q4_K Q4_0 Q5_K Q5_0 Q8_0 other
mmq_x_max 64 64 128 128 128 128 128 128 128
mmq_y 64 128 128 128 128 128 128 128 128
nwarps 4 8 8 8 8 8 8 8 8
A6 y1-prefetch yes no no no yes yes yes yes yes
MMQ_UNROLL=8 yes no no no no no no no no
N13 next-kb0 pf yes no no no no no no no no

Levers to optimize this kernel (what changed)

Commit 1 (ff4c006): Q6_K tuning

  • mmq_x_max=64, mmq_y=64, nwarps=4, MMQ_UNROLL=8 (stock WMMA 128/128/8/default)
  • load_tiles_q6_K hot path: pointer-stride qs walk, packed 6-bit-in-byte LDS staging
  • Y-pipe: prefetch second Y half into registers before load_tiles (A6 y1-prefetch)
  • Next-kb block-pointer prefetch (Q6_K-only)
  • Small-N decode: unroll k01 loop by 2 for mmq_x<=16
  • Q6_K epilogue dispatch on mmq_x: small=stock mul+add (bit-exact), large=fmaf (fast)

Commit 2 (9f07b0d): per-type geometry dispatch

  • Gate RDNA4 64/64/4 geometry on type==Q6_K; all others keep stock 128/128/8
  • A6 y1-prefetch gated on type != Q4_K (register pressure hurts Q4_K)

Commit 3 (fe2df7f): Q3_K + Q2_K fix

  • A6 gate tightened: type != Q4_K && type != Q3_K && type != Q2_K
    • Q3_K: scale-heavy K-quant, A6 register pressure regressed prefill -4.6% -> now +0.2%
    • Q2_K: A6 adds net overhead at both mmq_x=128 and 64, excluded
  • Q2_K mmq_x_max: 128 -> 64 on RDNA4
    • Stock 128 was suboptimal for 2-bit dequant-bound kernel (only ~2 TFLOPS)
    • mmq_x=64 gives 3x more blocks, smaller per-block work -> ~6.5 TFLOPS (+237%)

Requirements

  • I have read and agree with the contributing guidelines
  • AI usage disclosure: YES, used agent with kernel optimization skills guided , all changes reviewed and understood by the contributor. Commit messages marked with Assisted-by: Cursor.

qiongz and others added 3 commits July 12, 2026 17:57
@qiongz
MMQ: RDNA4 gfx1201 Q6_K prefill + small-N decode tuning
ff4c006 
Tune mul_mat_q for AMD Radeon AI PRO R9700 (gfx1201, RDNA4). All changes
are gated on RDNA4; non-RDNA4 targets compile to the upstream kernel.

Levers:
- Tile config: mmq_x_max=64, mmq_y=64, nwarps=4, MMQ_UNROLL=8 (stock WMMA
  128/128/8/default). Halves LDS -> 2 waves/CU, doubles dequant unroll.
- load_tiles_q6_K hot path: pointer-stride qs walk, optional first-row
  ql/qh prefetch, packed 6-bit-in-byte LDS staging expanded in vec_dot.
- Y-pipe: prefetch the second Y half into registers before load_tiles and
  copy it back for the second vec_dot, with the syncthreads required between
  the first vec_dot (reads tile_y) and the y1_reg->tile_y overwrite. Without
  that barrier cross-warp LDS races corrupt the dot product (abs_err up to
  80x, speculative-decoding accept collapses).
- Next-kb block-pointer prefetch (Q6_K).
- Small-N decode (mmq_x<=16): unroll the k01 loop by 2 so adjacent K-tile
  A-ldmatrix loads overlap the previous mma (ATT ds_read wait 6.5%->3.1%).
  Large mmq_x keeps the default unroll (unroll 2 raises VGPR, regresses
  prefill).
- Q6_K epilogue dispatch on mmq_x: small mmq_x<=16 uses the stock mul+add
  order (bit-identical to upstream, preserves speculative-decoding accept);
  large mmq_x uses hoisted-sd fmaf (fast, ~5e-5 fused-rounding residual
  harmless for prefill). fmaf's 1-rounding vs mul+add's 2 is the residual
  source; mul+add is bit-exact but ~32% slower on the prefill hot loop, so
  it is gated to the decode/verify path.

Micro-bench (mmq-bench, Q6_K, vs stock WMMA): ffn_gate N14 +12%, N16 +14%,
N24 +38%, N32 +35%; ffn_down N14 +23%, N32 +46%; attn_q N14 +5%; gdn N14
+22%; ffn_gate N512 +44%, N2048 +39%.

E2E (Qwen3.6-27B-UD-Q6_K_XL, llama-bench + server draft-mtp n-max=2):
pp512 705->876 (+24%), pp2048 687->853 (+24%), pp8192 636->775 (+22%);
tg128 20.8 (flat, N=1 uses MMVQ); MTP decode 45.7 t/s, accept 91%.

Precision: small-N (mmq_x<=16) bit-identical to upstream (ref_abs=0,
argmax 100%); large-N within 5e-5 (fmaf, prefill path).
@qiongz
MMQ: per-type geometry dispatch for RDNA4 (fix Q4/Q5/Q8 prefill)
9f07b0d 
The RDNA4 64/64/4 tile + A6 + UNROLL8 is optimal only for HBM-bound Q6_K
(low arithmetic intensity, dequant-heavy). Compute-bound types (Q4_K/Q4_0/
Q5_0/Q8_0, high AI) regressed -12..-23% because nwarps=4/mmq_y=64 cuts MMA
throughput. Gate the RDNA4-special geometry on type==Q6_K; all other types
fall through to stock 128/128/8. A6 y1-reg-prefetch is gated on
type != Q4_K (helps all prefill by avoiding a second HBM read of y1, but
Q4_K's scale-heavy dequant hits an occupancy cliff from y1_reg[] pressure
at mmq_x=128). N13 next-kb0 prefetch stays Q6_K-only.

mmq-bench (Qwen3.6-27B shapes, ffn_gate 17408x5120, vs stock, same-GPU A/B):
  Q6_K  N=512 +44.6%, N=2048 +41.4% (HBM-bound, keeps target)
  Q4_K  N=512 +0.2%,  N=2048 +0.1%
  Q4_0  N=512 +1.1%,  N=2048 +1.4%
  Q5_K  N=512 +6.3%,  N=2048 +6.8%
  Q5_0  N=512 +3.3%,  N=2048 +3.7%
  Q8_0  N=512 +1.2%,  N=2048 -0.1% (noise)
All types >= stock, no regression. Q6_K decode N=14 +9.2%, N=32 +37%.

Precision (vs stock golden, --check-ref): Q4/Q5/Q8 bit-exact (ref_abs=0,
ref_cos=1.0, argmax=1.0); Q6_K decode bit-exact, prefill 3.8e-5 fmaf
residual (accept-irrelevant, gated to mmq_x>16 by the epilogue dispatch).

E2E (Qwen3.6-27B-UD-Q6_K_XL, llama-server draft-mtp n-max=2, max_tokens=96
x3, temp=0, two-GPU A/B): accept 91.0% stock == 91.0% patched (identical);
t/s 45.5 == 45.6. MTP decode path (mmq_x=16) uses the stock mul+add epilogue
=> bit-exact => accept unchanged.

Bound-class (rocprof ATT): Q6_K wait=30.4%/WMMA=4.6% (HBM-bound, 64/64/4+A6
optimal); Q8_0 wait=17.2%/WMMA=11.8% (compute-bound, stock 128/128/8 optimal).

All gates are RDNA4-guarded; non-RDNA4 compiles to upstream stock. Only
mmq.cuh + mmq.cu touched (mmvq.cu/vecdotq.cuh N3 decode patch separate).
qiongz
MMQ: fix Q3_K prefill and Q2_K decode regressions on RDNA4
fe2df7f 
Tighten the A6 y1-reg-prefetch gate from type != Q4_K to also exclude
Q3_K and Q2_K, and cap Q2_K mmq_x_max at 64 (vs stock 128).

Q3_K (6-bit scales) hit the same occupancy cliff as Q4_K from y1_reg[]
pressure at mmq_x=128, regressing prefill -4.6%. Excluding it restores
prefill to stock (+0.2%).

Q2_K (2-bit, dequant-bound) had a pre-existing N=128 decode -70%
regression: mmq_x=128 gives only 136 blocks and A6 halves occupancy,
forcing a badly underutilized 2nd wave. mmq_x_max=64 doubles the block
count and alone lifts Q2_K prefill +237% (stock 128 was suboptimal for
the low-AI 2-bit kernel). A6 is also excluded - it adds net overhead at
both 128 and 64; the gain comes entirely from the tile-size change.

mmq-bench (ffn_gate 17408x5120, same-GPU0 A/B):
  Q2_K  N=512 +237%, N=2048 +214%; decode N=128 -0.2% (was -70%)
  Q3_K  N=512 +0.2%, N=2048 +1.1%; decode flat
  All other types unchanged (Q6_K +43%, Q5_K +6.6%, etc.)

Precision: bit-exact vs stock same-session. test-backend-ops 1134/1134
pass. All gates RDNA4-guarded; non-RDNA4 compiles to upstream stock.

Assisted-by: Cursor
@qiongz
qiongz requested a review from a team as a code owner July 12, 2026 12:42
@github-actions github-actions Bot added ggml changes relating to the ggml tensor library for machine learning CUDA Related to the CUDA backend labels Jul 12, 2026
@JohannesGaessler

JohannesGaessler commented Jul 12, 2026

Copy link
Copy Markdown
Contributor

Any arch-specific tuning needs to be rebased on top of #24127 .

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

No reviews

Assignees

No one assigned

Labels

CUDA Related to the CUDA backend ggml changes relating to the ggml tensor library for machine learning

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

2 participants

@qiongz @JohannesGaessler