q4_K mat-vec tuning for gfx1151 - #84
Draft
mgehre-amd wants to merge 1 commit into
Draft
Conversation
Port the structural choices ggml-vulkan's mul_mat_vecq.comp makes for q4_K into the
HIP mul_mat_vec_q path, driven by a side-by-side ISA comparison against the RADV
shader. On the gemma-4-31B-it Q4_K_M FFN gate/up shape the kernel goes from
292.5 us to 279.5 us, -4.4%, with no build flags.
End to end on that model, decode goes from 10.81 to 10.93 tok/s, +1.15% (95% CI
+0.68 to +1.63, 13 interleaved paired llama-bench runs, -p 128 -n 128 -d 128 -r 5
-fa 1 -mmp 0, 11/13 pairs positive). Prefill is the negative control and is
unchanged: it takes the MMQ path, which this does not touch. Measured -0.17%
median with a confidence interval spanning zero.
Everything is selected by defined(RDNA3_5) / GGML_CUDA_CC_IS_RDNA3_5, so NVIDIA,
CDNA, RDNA3.0 and RDNA4 compile byte-identical code to before.
* RADV-style vec_dot for q4_K plus a block_q8_1_x4 activation layout, worth 4.9%.
Each lane takes one aligned 16-byte chunk of qs and one nibble half, so the
weight fetch is a single 128-bit load instead of two 32-bit ones, and the
activation fetch is a single 128-bit load of 16 contiguous bytes. Two lanes
share a chunk and read the same 16 bytes with different shifts, the same
register-level redundancy RADV accepts in exchange for wide loads. Still 16
threads per superblock, so VDR stays 2 and the K-loop trip count is unchanged.
The two halves are a package: the x4 layout without the wide-load vec_dot is a
6.6% regression, worse than reverting both.
* A 4x manual unroll of the q4_K K-loop with a sched_group_barrier, worth 1.9%.
The trip count is runtime-dependent so the compiler will not unroll it. The
barrier is load-bearing rather than a tuning knob: unrolling without it costs
10.4%, far worse than not unrolling at all, because the whole unrolled body's
live set is scheduled at once. The original loop is preserved verbatim for
every other type and architecture.
* nwarps 2 instead of 1 for q4_K on the RDNA2 parameter table, worth 1.0%. One
wave per output row leaves only two q4_K superblocks in flight.
Note the x4 layout stores d*sum(q) in ds.y, not sum(x) as the plain block_q8_1
layout does. This mirrors ggml-vulkan's quantize_q8_1 and is what the q4_K min term
consumes when it reads the block sum instead of recomputing it with 4 extra dp4a.
The two differ by the quantization residual, and using sum(x) there raises the error
on this shape from 25e-6 to a data-dependent 35-216e-6 against a 5e-4 tolerance -
silent, and it only trips the test about once in 30 runs. The convention is safe to
define here because the x4 layout is gated to q4_K on RDNA3.5 and has exactly one
reader.
Also measured and rejected: a wide dm+scales head load, wide qs loads, dword scale
loads, CU mode, rows-per-block 2/4, unroll 2/5/6/8, and five mechanisms for lowering
VGPR (sched_barrier, a second sched_group_barrier pattern, iglp_opt,
__launch_bounds__ min-blocks, amdgpu_waves_per_eu). VGPR stays at 89-93 with zero
spills under all of them versus RADV's 48; register pressure did not track
performance in this kernel.
Wave64 for the mat-vec kernel is worth a further 1.6% but is left to a follow-up:
-mwavefrontsize64 is a per-TU flag and flips all 265 kernels in mmvq.cu, not just
the q4_K one, and the other 264 are unmeasured. Scoping it needs the q4_K
instantiation moved to its own translation unit. A function attribute cannot
substitute - clang refuses to inline any wave32 function into a
target("wavefrontsize64") one, down to the threadIdx accessors.
Measurement: rocprofv3 --kernel-trace mean over 3465 dispatches of
MUL_MAT type_a=q4_K m=21504 n=1 k=5376 with GGML_CUDA_DQ_MMV=0, four interleaved
passes against the base under an exclusive GPU lock at DPM high, gated on the GPU
edge sensor at 52 C. GGML_CUDA_DQ_MMV=0 is required to reach mul_mat_vec_q at all
for q4_K here: at ne11 == 1 with k % QK_K == 0 the dispatcher early-returns into
mul_mat_vec_dq_q4_K. In the model the op is reached anyway, because it is fused
(MUL_MAT + MUL_MAT + GLU) and gemma-4 uses GEGLU, which falls through the
SWIGLU-only dq_glu guard.
Accuracy, 20 independent draws against the CPU reference, tolerance 5e-4: base
median 24e-6, this change 25e-6, ggml-vulkan/RADV 25e-6 - indistinguishable.
Verification: test-backend-ops test -o
MUL_MAT,MUL_MAT_ID,MUL_MAT_VEC_FUSION,MUL_MAT_ID_FUSION passes with and without
GGML_CUDA_DQ_MMV=0, covering the fused instantiation the model dispatches, which
-o MUL_MAT alone does not match.
Assisted-by: Claude Opus 5
mgehre-amd
force-pushed
the
matthias.q4k-mmvq-gfx1151
branch
from
August 13, 2026 17:12
e18b2e6 to
1b600a3
Compare
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Ports the structural choices ggml-vulkan's
mul_mat_vecq.compmakes for q4_K into the HIPmul_mat_vec_qpath, driven by a side-by-side ISA comparison against the RADV shader.Two commits: the first is the tuning work, the second reduces it to the parts that measurably
contribute and turns them on by architecture instead of by build flag.
Result
Shape:
MUL_MAT type_a=q4_K type_b=f32 m=21504 n=1 k=5376- the FFN gate/up matvec, the largestsingle op in gemma-4-31B-it Q4_K_M decode (60 dispatches/token, 39% of decode in a per-op profile).
Kernel time,
rocprofv3 --kernel-tracemean over 3465 dispatches, gfx1151 / Radeon 8060S,ROCm 7.15.0a20260728, median of 4 interleaved passes:
gfx11baseFor reference RADV's
mul_mat_vec_q4_k_q8_1_f32measures ~280 us on the same shape.What changed
Selected by
defined(RDNA3_5)/GGML_CUDA_CC_IS_RDNA3_5, so no other architecture changesbehaviour and no macros need to be passed:
vec_dot+block_q8_1_x4activation layout (4.9%). Each thread takes onealigned 16-byte chunk of
qsand one nibble half, giving a 128-bit weight load, plus 16contiguous activation bytes for a 128-bit activation load. These are a package: the x4 layout
without the wide-load
vec_dotis a 6.6% regression.sched_group_barrier(1.9%), scoped to q4_K. The barrier isload-bearing - unrolling without it costs 10.4%, far worse than not unrolling.
nwarpsentry for q4_K, set to 2 (1.0%).No build flags are required - the win is in the default build.
Deliberately not here: wave64 for the mat-vec kernel, worth a further 1.6%. It is split to a
follow-up because
-mwavefrontsize64is a per-TU flag that flips all 265 kernels inmmvq.cu,not just the q4_K one, and the other 264 are unmeasured. Scoping it needs the q4_K instantiation
moved to its own translation unit. A function attribute cannot substitute: clang refuses to
inline any wave32 function into a
target("wavefrontsize64")one, down tothreadIdxaccessors.Accuracy
The x4 layout stores
d*sum(q)inds.y, notsum(x)as the plainblock_q8_1layout does.This matches ggml-vulkan's
quantize_q8_1and is the quantity the q4_K min term consumes when itreads the block sum instead of recomputing it with four extra
dp4a.Getting this wrong is silent and data-dependent, so it is worth stating: an earlier revision of
this branch fed
sum(x)to that min term, which raised the error on this shape from 25e-6 to adata-dependent 35-216e-6 against a 5e-4 tolerance, and tripped the test roughly once in 30 runs.
Error against the CPU reference, 20 independent draws (tolerance 5e-4):
gfx11baseAll three are indistinguishable, so the change is accuracy-neutral against base and against RADV.
Verification
test-backend-ops test -o MUL_MAT,MUL_MAT_ID,MUL_MAT_VEC_FUSION,MUL_MAT_ID_FUSION -b ROCm0passes with and without
GGML_CUDA_DQ_MMV=0. This covers thefused instantiation the model actually dispatches (gemma-4 is GEGLU, so it falls through the
SWIGLU-only
dq_gluguard intomul_mat_vec_qwithhas_fusion=true);-o MUL_MATalone doesnot match it.
GGML_CUDA_DQ_MMV=0is required to reachmul_mat_vec_qat all for q4_K on gfx1151: atne11 == 1withk % QK_K == 0the dispatcher early-returns intomul_mat_vec_dq_q4_K.Reproducing
Runs were serialised under exclusive GPU access at DPM level high and gated on the GPU edge
sensor at 52 C; long unsegmented batches were shown to throttle and degrade results by up to 24%.