HIP: fix MMQ tile-barrier race and tune tile width for RDNA3.5

[roberteg16]

Summary

Two commits: a correctness fix that stands alone, and a tile-width change measured on top of it.

1. Missing tile barrier at the top of the MMQ K loop. load_tiles() overwrites tile_x while other warps of the same block can still be inside the previous iteration's vec_dot(). The guard there was mmq_hip_tile_barrier<mmq_x>(), which is elided for mmq_x <= 32, and the __syncthreads() calls further down all sit between the y tile store and vec_dot, on the wrong side of the hazard.

With nwarps=4 on RDNA3.5 that left mmq_x of 16 and 32 unprotected for the whole K loop. 115 of 1568 measured configurations returned wrong results, non-deterministically — the same q8_0 8192x2048 case repeated three times gave max abs errors of 9.49, 1.37 and 4.44 against a correct value of 3e-05 — and test-backend-ops -o MUL_MAT_ID failed intermittently on q4_0. MoE is worst hit, since every MoE launch on this architecture runs at mmq_x=16. After the fix all 1568 configurations agree and repeat runs are bit-identical. Cost is 1.3% at mmq_x=16, nothing above 32.

2. Tuned tile width for RDNA3.5. Sweeping the 99 shapes that 20 models dispatch to mul_mat_q, best width per shape vs. what the heuristic picks:

tokens per ubatch	dense	MoE
512	90 shapes, 1.004x	9 shapes, 1.172x
2048	88 shapes, 1.000x	9 shapes, 1.485x

Dense is already optimal and is left alone, except for mmq_x=64 with q8_0 — a 6-7x cliff the heuristic hits whenever ncols_to_tile falls in (48, 64], i.e. a 64-token batch. MoE gets a 32-wide tile instead: the compacted per-expert grid already supplies the parallelism that column tiling gives a dense GEMM.

Measured end to end

Alternated builds, -ub 512 unless noted. Whether the change fires depends on
ncols_to_tile = min(2*ceil(M*top_k/n_experts), M), so it moves with both the batch size and the routing shape.

batch	Qwen3.5/3.6-35B-A3B	gemma-4-26B-A4B	GLM-4.7-Flash
pp64	1.74x	not measured	1.21-1.31x
pp128	neutral	neutral	neutral
pp512	neutral	+11.8%	+8.5%
pp2048	neutral	+15.1%	+8.6%
pp2048 -ub 2048	+13.0-13.8%	+9.4%	+3.4%

pp128 is neutral by construction: ncols_to_tile lands on 8-16 there, which the heuristic already picks, so the rule is a no-op and only the barrier's 1.3% remains — below the noise floor.

The -ub 2048 row comes from a 20-model run whose 16 dense models go through identical code on this path and act as a negative control: median 0.992x, range 0.963-1.049x.

Test plan

• test-backend-ops test -o MUL_MAT -b ROCm0 and -o MUL_MAT_ID -b ROCm0 — pass, and MUL_MAT_ID no longer fails intermittently
• llama-bench pp64/128/512/2048 at the default -ub, and pp2048 -ub 2048 across 20 models (0.5B-35B, dense and MoE), builds alternated
• Selected width confirmed with GGML_CUDA_MMQ_DEBUG=1 before each A/B, so a build that does not actually apply the rule cannot be mistaken for a null result
• Both commits build independently (bisectable)
• Review whether the unconditional __syncthreads() should be restricted to RDNA3.5: it is active on every backend and the same hazard applies there, but it has only been validated on gfx1151

Notes

• Base is gfx11, not master.
• Measurements come from the harness in #, which is independent and can land in either order.

[mgehre-amd]

Thanks!

[Copilot]

🟢 Ready to approve

The synchronization fix addresses a concrete race, and the RDNA3.5 tuning is scoped and guarded by existing granularity/shared-memory eligibility checks.

This review doesn't count toward merge requirements. Sign up for the private preview to control whether Copilot approvals count.

Pull request overview

Fixes a correctness race in the HIP MMQ kernel’s K-loop (shared tile_x reuse across iterations) and adjusts RDNA3.5-specific tile-width selection to avoid known slow configurations and improve MoE throughput.

Changes:

• Add a block-wide synchronization at the top of the MMQ K-loop before load_tiles() to prevent tile_x being overwritten while other warps are still executing the prior iteration’s vec_dot().
• Introduce an RDNA3.5 HIP-only tuned override for mmq_x (notably avoiding the mmq_x=64 Q8_0 cliff, and preferring narrower tiles for MoE), applied only when the tuned width is valid for the shape (granularity + shared-memory constraints).

File summaries

File	Description
ggml/src/ggml-cuda/mmq.cuh	Adds a missing synchronization in the MMQ K-loop and applies RDNA3.5 HIP-only mmq_x tuning with validity guards.

Review details

• Files reviewed: 1/1 changed files
• Comments generated: 0
• Review effort level: Low

We're testing this review assessment. Please use 👍 or 👎 to tell us if it's correct.