Skip to content

q4_K mat-vec tuning for gfx1151 - #84

Draft
mgehre-amd wants to merge 1 commit into
gfx11from
matthias.q4k-mmvq-gfx1151
Draft

q4_K mat-vec tuning for gfx1151#84
mgehre-amd wants to merge 1 commit into
gfx11from
matthias.q4k-mmvq-gfx1151

Conversation

@mgehre-amd

Copy link
Copy Markdown
Collaborator

Ports 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.

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 largest
single 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-trace mean over 3465 dispatches, gfx1151 / Radeon 8060S,
ROCm 7.15.0a20260728, median of 4 interleaved passes:

Build Kernel time vs base
gfx11 base 292.5 us -
This PR (no build flags) 279.5 us -4.4%

For reference RADV's mul_mat_vec_q4_k_q8_1_f32 measures ~280 us on the same shape.

What changed

Selected by defined(RDNA3_5) / GGML_CUDA_CC_IS_RDNA3_5, so no other architecture changes
behaviour and no macros need to be passed:

  • RADV-style q4_K vec_dot + block_q8_1_x4 activation layout (4.9%). Each thread takes one
    aligned 16-byte chunk of qs and one nibble half, giving a 128-bit weight load, plus 16
    contiguous activation bytes for a 128-bit activation load. These are a package: the x4 layout
    without the wide-load vec_dot is a 6.6% regression.
  • 4x K-loop unroll with sched_group_barrier (1.9%), scoped to q4_K. The barrier is
    load-bearing - unrolling without it costs 10.4%, far worse than not unrolling.
  • RDNA2-table nwarps entry 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 -mwavefrontsize64 is a per-TU flag that 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 threadIdx accessors.

Accuracy

The x4 layout stores d*sum(q) in ds.y, not sum(x) as the plain block_q8_1 layout does.
This matches ggml-vulkan's quantize_q8_1 and is the quantity the q4_K min term consumes when it
reads 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 a
data-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):

Backend min median max
gfx11 base 0.000023 0.000024 0.000026
This PR 0.000023 0.000025 0.000027
Vulkan / RADV 0.000024 0.000025 0.000026

All 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 ROCm0
passes with and without GGML_CUDA_DQ_MMV=0. This covers the
fused instantiation the model actually dispatches (gemma-4 is GEGLU, so it falls through the
SWIGLU-only dq_glu guard into mul_mat_vec_q with has_fusion=true); -o MUL_MAT alone does
not match it.

GGML_CUDA_DQ_MMV=0 is required to reach mul_mat_vec_q at all for q4_K on gfx1151: at
ne11 == 1 with k % QK_K == 0 the dispatcher early-returns into mul_mat_vec_dq_q4_K.

Reproducing

cmake -B build -DGGML_HIP=ON -DAMDGPU_TARGETS=gfx1151 -DCMAKE_BUILD_TYPE=Release
cmake --build build --target test-backend-ops -j$(nproc)

GGML_CUDA_DQ_MMV=0 ./build/bin/test-backend-ops perf -o MUL_MAT -b ROCm0 \
    -p 'type_a=q4_K,type_b=f32,m=21504,n=1,k=5376'

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%.

@mgehre-amd mgehre-amd changed the title # q4_K mat-vec tuning for gfx1151 q4_K mat-vec tuning for gfx1151 Aug 13, 2026
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
mgehre-amd force-pushed the matthias.q4k-mmvq-gfx1151 branch from e18b2e6 to 1b600a3 Compare August 13, 2026 17:12
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant