Skip to content

[2/2] cuda: TQ4_1S MMQ prefill on RDNA3 and RDNA2 - #337

Merged
TheTom merged 6 commits into
TheTom:feature/turboquant-kv-cachefrom
jasstrong:pr/tq4_1s-mmq-rdna
Sep 3, 2026
Merged

[2/2] cuda: TQ4_1S MMQ prefill on RDNA3 and RDNA2#337
TheTom merged 6 commits into
TheTom:feature/turboquant-kv-cachefrom
jasstrong:pr/tq4_1s-mmq-rdna

Conversation

@jasstrong

Copy link
Copy Markdown

Series: TQ4_1S MMQ prefill on AMD — part 2/2

Depends on #336 and is branched on it, so GitHub shows its commit here until #336 lands. This part is not severable: it configures and gates a path that #336 introduces.

Overview

Extends TQ4_1S MMQ prefill from CDNA to RDNA3 and RDNA2. Very little was actually missing — the vec_dot is the generic q8_0_16 one already shared with Q3_K, and MMQ already dispatches WMMA on RDNA3 by itself. Three things stood in the way.

1. The gate. The entry points tested amd_mfma_available(cc), which is CDNA-only. Widened to accept WMMA and RDNA2.

2. Missing tile configs. TQ4_1S CASE entries existed only in mmq-config-cdna.cuh, so ggml_cuda_mmq_get_config_rdna3() returned GGML_TYPE_COUNT for every J and the J-selection loop aborted with J_best=0. Added entries mirroring Q3_K, whose shared-memory layout TQ4_1S already borrows.

3. A real bug in the tile loader, which is what RDNA2 needed. ggml_cuda_mmq_load_tiles_tq4_1s wrote a single hardcoded tile layout and carried the comment "TQ4_1S MMQ is MFMA-only". Every other loader branches on AMD_MFMA_AVAILABLE / TURING_MMA_AVAILABLE / AMD_WMMA_AVAILABLE, because dp4a consumes a different shared-memory layout: rows padded by one int, and the scales based after that padded quant region rather than at a fixed 2*MMQ_TILE_NE_K offset. mmq_get_dp4a_tile_x_sizes already mapped TQ4_1S to MMQ_DP4A_TXS_Q8_0_16, so the sizing was right all along — only the loader ignored it, writing quants and scales where the vec_dot never read them. TQ4_1S was also missing from the dp4a traits block and so picked up the wrong vec_dot.

Additional information

Qwen3.6-35B-A3B-TQ4max, ROCm 7.2.3, each card measured against itself with the path disabled:

GPU arch prefill off → on speedup ppl off → on tq4_1s MUL_MAT
7900XTX gfx1100 (WMMA) 759.0 → 2533.2 3.34x 3.2698 → 3.2595 269/269
V620 gfx1030 (dp4a, no matrix cores) 424.0 → 1983.7 4.68x 3.2705 → 3.2660 269/269

Decode is unchanged on both. The V620 number is the interesting one: 4.68x prefill on hardware with no matrix cores at all, reaching about 86% of an MI210's prefill.

The two commits are split so each is independently bisectable — RDNA3 (configs + WMMA gate) does not depend on the loader work; RDNA2 (loader + traits + configs + gate) does.

GGML_TQ_LPR was also swept 2→32 on both cards and is flat within ±0.2, so the wave64-tuned default of 16 is safe on wave32 and no per-arch tuning is needed.

Requirements

  • I have read and agree with the contributing guidelines
  • AI usage disclosure: YES — Claude Code (Opus 5) was used: it diagnosed the J_best=0 abort and the dp4a layout bug, wrote the fix and commit messages, and ran the correctness and perplexity verification on the 7900XTX and V620. I reviewed the change and am responsible for every line of it.

jas and others added 3 commits August 31, 2026 21:59
TQ4_1S weights previously had no MMQ path, so prefill dequantised them to
f16 and went through hipBLAS. That threw away the format's whole point:
the 5bpw footprint was paid for at load time and then not used at compute
time.

Add a TQ4_1S tile loader, MMQ type traits, and dispatch, and pre-rotate
the activation so the block-local turbo WHT cancels against the weights.
The weights stay in their rotated-domain int8 centroid form, so the stock
MFMA-i8 MMQ kernel consumes them directly. The type reuses the Q3_K
shared-memory layout, since its per-16 scales have the same shape, and so
takes the same tile geometry Q3_K and IQ2_XS use on CDNA (occupancy 1,
I=128); other geometries do not match what that layout expects and write
out of bounds.

Measured on an MI210 with Qwen3.6-35B-A3B-TQ4max, against this same
branch with the path disabled:

  prefill  971 -> 2124 t/s  (2.19x)
  decode    95 ->   93 t/s  (-2.6%, this is a prefill-only path)
  ppl     3.2653 -> 3.2653  (identical)

It also keeps roughly 13GB less VRAM than converting the experts to Q8_0.
All 147 tq4_1s MUL_MAT cases pass.

The path is env-gated behind GGML_TQ_MMQ so it can be A/B'd against the
hipBLAS fallback.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01NxP6x5bmUDYFvmouceN2mR
The MMQ entry points were gated on amd_mfma_available(), which is CDNA
only, so RDNA3 fell back to dequantising to f16 and calling hipBLAS.
Nothing in the kernel actually needs MFMA: the vec_dot is the generic
q8_0_16 one already shared with Q3_K, and upstream MMQ dispatches WMMA
on RDNA3 by itself. The only missing piece was the per-arch tile
configuration, which existed solely in mmq-config-cdna.cuh.

Mirror Q3_K's twelve RDNA3 entries for TQ4_1S, which is already the
shared-memory layout it borrows, and let the gate accept WMMA.

7900XTX: prefill 759 -> 2533 t/s (3.34x), decode unchanged, 269 of 269
tq4_1s cases pass, perplexity 3.2595 with MMQ against 3.2698 without.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01NxP6x5bmUDYFvmouceN2mR
ggml_cuda_mmq_load_tiles_tq4_1s wrote a single hardcoded tile layout and
was commented "TQ4_1S MMQ is MFMA-only". Every other loader branches on
AMD_MFMA_AVAILABLE / TURING_MMA_AVAILABLE / AMD_WMMA_AVAILABLE, because
dp4a consumes a different shared-memory layout: rows are padded by one
int, and the scales start after that padded quant region rather than at
a fixed 2*MMQ_TILE_NE_K offset.

mmq_get_dp4a_tile_x_sizes already mapped TQ4_1S to MMQ_DP4A_TXS_Q8_0_16,
so the sizing was right all along; only the loader ignored it, writing
quants and scales where the vec_dot never read them. TQ4_1S was also
missing from the dp4a traits block, so it picked up the wrong vec_dot.

With both fixed, TQ4_1S MMQ runs on hardware with no matrix cores at all.
V620 (gfx1030): prefill 424 -> 1984 t/s (4.68x), decode unchanged, 269 of
269 tq4_1s cases pass, perplexity 3.2660 with MMQ against 3.2705 without.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01NxP6x5bmUDYFvmouceN2mR
…and MUSA

tq4_1s_qs_to_int8x4() called the AMD-only byte-permute builtin unconditionally,
so the CUDA and MUSA CI builds failed to compile mmq-load-tiles.cuh. Keep
v_perm_b32 on HIP and use the equivalent __byte_perm selector elsewhere. The
TQ4_1S MMQ path itself is still gated to AMD at dispatch time; this only makes
the header compile on the other toolchains.
@TheTom

TheTom commented Sep 3, 2026

Copy link
Copy Markdown
Owner

Reviewed. Your two commits are clean: the dp4a tile layout matches ggml_cuda_mmq_load_tiles_iq2_xs line for line (same 2*MMQ_TILE_NE_K + 1 quant stride, same scale base), the d0/d1 slot placement matches the per-4-int granularity q8_0_16 expects, and the RDNA3 configs mirror Q3_K as described.

The only build failure is the __builtin_amdgcn_perm inherited from #336, which is now fixed on that branch (03ee3ed). The Windows x64-vulkan failure is test-thread-safety dying with 0xc0000374 (heap corruption); it hit #336 as well but not #338, which shares the same code, so I'm treating it as a runner flake for now and will watch the rerun.

Plan: land #336 when its CI is green, then push the guard commit here so this PR re-runs, then merge. One follow-up for either PR: getenv("GGML_TQ_MMQ") runs per mul_mat call. Short-circuited away on NVIDIA, but on AMD it is a libc call per node. A static const bool like the other knobs would be better.

@TheTom

TheTom commented Sep 3, 2026

Copy link
Copy Markdown
Owner

#336 is merged. Pushed a merge of its perm guard (03ee3ed) here so CI re-runs with the fix; will merge once cuda/musa/hip are green.

getenv() ran on every mul_mat dispatch. It is short-circuited away on
NVIDIA but is a libc call per node on the AMD path. Read it once through
a static, like the other knobs.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01NxP6x5bmUDYFvmouceN2mR
@TheTom
TheTom merged commit ed0816b into TheTom:feature/turboquant-kv-cache Sep 3, 2026
6 of 21 checks passed
@jasstrong

Copy link
Copy Markdown
Author

Heads-up: the "read GGML_TQ_MMQ once" change I pushed here (439fe67) initializes its static from a call to itself, so on AMD the first TQ mul_mat with n >= 8 deadlocks on the static's guard. Every prefill hangs with GGML_TQ_MMQ=1 on RDNA and CDNA; NVIDIA never calls the function. One-line fix with MI210 results in #348. Sorry for the trouble.

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

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants