Skip to content

HIP : RDNA3.5 MMVQ/MMQ kernels and per-type dispatch thresholds for quantised matmul #56

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
SimonTeixidor wants to merge 15 commits into
base: master
Choose a base branch
from
Open

HIP : RDNA3.5 MMVQ/MMQ kernels and per-type dispatch thresholds for quantised matmul #56

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
15 commits
Select commit Hold shift + click to select a range
b358475
CUDA/HIP: let IQ2_XS/IQ2_S reach hipBLAS above ne11=128 on RDNA3.5
SimonTeixidor Sep 7, 2026
2ae2ee9
HIP: branch-free SWAR for __vsub4 / __vcmpne4 / __vcmpeq4
SimonTeixidor Sep 8, 2026
66efda4
CUDA/HIP: avoid emulated __vcmpne4/__vsub4 in the IQ2/IQ3 sign path
SimonTeixidor Sep 7, 2026
e1430b1
CUDA/HIP: route the fork-only IQ3_S sign sites through apply_signs4
SimonTeixidor Sep 8, 2026
df3e4d7
HIP: per-type MMVQ batch thresholds on RDNA3.5
SimonTeixidor Sep 8, 2026
856d62f
HIP: two MMVQ rows per block on RDNA3.5 for ncols_dst 2..8
SimonTeixidor Sep 9, 2026
b9bbaa7
HIP: prefer the vector kernels below one MMQ tile on RDNA3.5
SimonTeixidor Sep 9, 2026
c6e682c
HIP: full-rate 24-bit multiply for the IQ2_XS/IQ2_S/IQ3_S scale epilogue
SimonTeixidor Sep 9, 2026
9a08368
HIP: two output rows per block in the RDNA3.5 four-column MMVQ kernels
SimonTeixidor Sep 9, 2026
dada19e
CUDA/HIP: make calc_rows_per_block a per-type table on RDNA3.5
SimonTeixidor Sep 9, 2026
7ed89d9
CUDA/HIP: cheaper MMVQ epilogue and sign application for the codebook…
SimonTeixidor Sep 10, 2026
c4d485b
HIP: env override for the RDNA3.5 per-type MMVQ/MMQ crossover
SimonTeixidor Sep 10, 2026
151cdf1
hip: extend the RDNA3.5 MMQ prefetch to the J = 16 tile for Q4_K/Q5_K…
SimonTeixidor Sep 10, 2026
16730c2
HIP: re-tune the whole RDNA3.5 MMVQ/MMQ crossover table
SimonTeixidor Sep 11, 2026
79194d3
HIP: extend the RDNA3.5 crossover re-tune to the eleven unmeasured types
SimonTeixidor Sep 11, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion ggml/src/ggml-cuda/ggml-cuda.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1859,7 +1859,7 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor
ggml_cuda_mul_mat_f(ctx, src0, src1, nullptr, dst);
return;
}
if (ggml_cuda_should_use_mmvq(src0->type, cc, ne11)) {
if (ggml_cuda_should_use_mmvq(src0->type, cc, ne01, ne11, hint == GGML_HINT_EXACT_BATCH)) {
ggml_cuda_mul_mat_vec_q(ctx, src0, src1, nullptr, dst);
return;
}
Expand Down
41 changes: 13 additions & 28 deletions ggml/src/ggml-cuda/mmq-load-tiles.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1228,13 +1228,10 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
#pragma unroll
for (int l = 0; l < QR2_XXS; ++l) {
const uint2 grid_pos = ((const uint2*)iq2xxs_grid)[aux8[l]];
const uint32_t signs = unpack_ksigns(aux32 >> (7 * l));
const uint32_t signs = unpack_ksigns8(aux32 >> (7 * l));

const int signs0 = __vcmpne4(signs & 0x08040201, 0);
const int grid0 = __vsub4(grid_pos.x ^ signs0, signs0);

const int signs1 = __vcmpne4(signs & 0x80402010, 0);
const int grid1 = __vsub4(grid_pos.y ^ signs1, signs1);
const int grid0 = apply_signs4(grid_pos.x, signs);
const int grid1 = apply_signs4(grid_pos.y, signs >> 4);

#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
x_qs[i*sram_stride + 8*kqsx + (2*l + 0)] = grid0;
Expand Down Expand Up @@ -1291,13 +1288,10 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
#pragma unroll
for (int l = 0; l < QR2_XS; ++l) {
const uint2 grid_pos = ((const uint2*)iq2xs_grid)[q2[l] & 0x1FF];
const uint32_t signs = unpack_ksigns(q2[l] >> 9);

const int signs0 = __vcmpne4(signs & 0x08040201, 0);
const int grid_l = __vsub4(grid_pos.x ^ signs0, signs0);
const uint32_t signs = unpack_ksigns8(q2[l] >> 9);

const int signs1 = __vcmpne4(signs & 0x80402010, 0);
const int grid_h = __vsub4(grid_pos.y ^ signs1, signs1);
const int grid_l = apply_signs4(grid_pos.x, signs);
const int grid_h = apply_signs4(grid_pos.y, signs >> 4);

#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
x_qs[i*sram_stride + 8*kqsx + (2*l + 0)] = grid_l;
Expand Down Expand Up @@ -1361,11 +1355,8 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
for (int l = 0; l < QR2_S; ++l) {
const int * grid_pos = (const int *)(iq2s_grid + (qs[l] | ((qh << (8-2*l)) & 0x300)));

const int signs0 = __vcmpne4(((signs_packed_8[l] & 0x03) << 7) | ((signs_packed_8[l] & 0x0C) << 21), 0x00000000);
const int signs1 = __vcmpne4(((signs_packed_8[l] & 0x30) << 3) | ((signs_packed_8[l] & 0xC0) << 17), 0x00000000);

const int grid_l = __vsub4(grid_pos[0] ^ signs0, signs0);
const int grid_h = __vsub4(grid_pos[1] ^ signs1, signs1);
const int grid_l = apply_signs4(grid_pos[0], signs_packed_8[l]);
const int grid_h = apply_signs4(grid_pos[1], signs_packed_8[l] >> 4);

#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
x_qs[i*sram_stride + 8*kqsx + (2*l + 0)] = grid_l;
Expand Down Expand Up @@ -1425,13 +1416,10 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
#pragma unroll
for (int l = 0; l < QR3_XXS; ++l) {
const int2 grid_pos = make_int2(iq3xxs_grid[q3[2*l+0]], iq3xxs_grid[q3[2*l+1]]);
const uint32_t signs = unpack_ksigns(aux32 >> (7*l));
const uint32_t signs = unpack_ksigns8(aux32 >> (7*l));

const int signs0 = __vcmpne4(signs & 0x08040201, 0);
const int grid_l = __vsub4(grid_pos.x ^ signs0, signs0);

const int signs1 = __vcmpne4(signs & 0x80402010, 0);
const int grid_h = __vsub4(grid_pos.y ^ signs1, signs1);
const int grid_l = apply_signs4(grid_pos.x, signs);
const int grid_h = apply_signs4(grid_pos.y, signs >> 4);

#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
x_qs[i*sram_stride + 8*kqsx + (2*l + 0)] = grid_l;
Expand Down Expand Up @@ -1496,11 +1484,8 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
iq3s_grid[qs[2*l+0] | ((qh << (8 - 2*l)) & 0x100)],
iq3s_grid[qs[2*l+1] | ((qh << (7 - 2*l)) & 0x100)]);

const int signs0 = __vcmpne4(((signs_packed_8[l] & 0x03) << 7) | ((signs_packed_8[l] & 0x0C) << 21), 0x00000000);
const int signs1 = __vcmpne4(((signs_packed_8[l] & 0x30) << 3) | ((signs_packed_8[l] & 0xC0) << 17), 0x00000000);

const int grid_l = __vsub4(grid_pos.x ^ signs0, signs0);
const int grid_h = __vsub4(grid_pos.y ^ signs1, signs1);
const int grid_l = apply_signs4(grid_pos.x, signs_packed_8[l]);
const int grid_h = apply_signs4(grid_pos.y, signs_packed_8[l] >> 4);

#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
x_qs[i*sram_stride + 8*kqsx + (2*l+0)] = grid_l;
Expand Down
2 changes: 1 addition & 1 deletion ggml/src/ggml-cuda/mmq.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -563,7 +563,7 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t
return ne11 <= (GGML_CUDA_CC_IS_RDNA3_0(cc) ? 128 : 256);
case GGML_TYPE_IQ2_XS:
case GGML_TYPE_IQ2_S:
return GGML_CUDA_CC_IS_RDNA3_5(cc) || ne11 <= 128;
return ne11 <= 128;
default:
return true;
}
Expand Down
36 changes: 33 additions & 3 deletions ggml/src/ggml-cuda/mmq.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -876,10 +876,40 @@ template <ggml_type type, int J, bool fallback>
static constexpr __host__ __device__ bool ggml_cuda_mmq_use_prefetch() {
#if defined(RDNA3_5) && defined(AMD_WMMA_AVAILABLE)
// Whitelist: the extra registers cause spills in several other specializations.
//
// The J = 16 entries were added on 2026-09-10 (results/2026-09-10-mmq-j16-prefetch/). J = 16 is
// the only tile MMQ instantiates for ne11 <= 16 on this arch, i.e. all of single-stream
// decode, small-batch decode and speculative verification; the original sweep was
// MUL_MAT_ID at 2048 tokens and never reached it. Q4_K J = 16 and the already-whitelisted
// Q4_K J = 48 both sit at 2 blocks resident per 64 KiB CU (22080 / 26816 B of LDS), so the
// "only 1-2 blocks are resident" argument above applies unchanged at J = 16, while the y
// staging costs 10 VGPRs there against 36 at J = 128.
//
// TAKEN, per call at ne11 = 6 (m = 17408, k = 5120), 6 arms per side, balanced order:
// Q4_K J=16 -3.43 % (t = -6.18, CI -4.69 .. -2.16), and -2.5 .. -3.7 % across the whole
// ne11 = 3..16 range that dispatches this tile, null at ne11 = 1,2 (MMVQ).
// Q5_K J=16 -1.1 % typical, and Q6_K J=16 -0.3 %, neither resolved on its own against a
// +/-1.3 % control band (untouched q2_K/iq2_s cases moved that much). Taken on
// the group end-to-end result plus a +7 / +12 VGPR cost, not on their own
// per-call numbers; they are the weak half of this entry.
// End-to-end decode at 6 sequences, the three cells above being the only changed kernels at
// that batch: +0.80 % (CI +0.18 .. +1.42) on Qwen3.8-27B-UD-IQ4_XS where they are 26.0 % of
// the weights, and +4.65 % (CI +3.51 .. +5.78) on Qwen3.8-27B-Q4_K_M where they are 93.6 %.
// 1 sequence and prefill are null on both. Outputs stay bit-identical (equal per-chunk
// perplexity at -ub 16, against a base-vs-base determinism control).
//
// REFUSED, and why, so the measurement that excludes them is on the record:
// IQ4_XS J=16, IQ3_S J=16 no per-call effect at any ne11 = 7..16 that reaches this tile
// (every CI crosses zero), and they are the expensive cells: +33 .. +50 VGPR,
// which puts IQ4_XS fallback at 247 (mul_mat_q) and 251 (routed_compact) of
// 256. Cost with no measured benefit; not taken.
// Q2_K J=16, Q3_K J=16, IQ2_S J=16 already at 256 VGPR and already spilling on this
// branch without any prefetch (51 / 3 / 7 registers), so staging more would
// only deepen the spill. Not measured, not taken.
return (type == GGML_TYPE_Q8_0 && (J == 48 || J == 128) && !fallback) ||
(type == GGML_TYPE_Q6_K && J == 32) ||
(type == GGML_TYPE_Q5_K && J == 32) ||
(type == GGML_TYPE_Q4_K && J == 48) ||
(type == GGML_TYPE_Q6_K && (J == 16 || J == 32)) ||
(type == GGML_TYPE_Q5_K && (J == 16 || J == 32)) ||
(type == GGML_TYPE_Q4_K && (J == 16 || J == 48)) ||
(type == GGML_TYPE_IQ2_S && J == 128) ||
(type == GGML_TYPE_IQ3_XXS && J == 128);
#else
Expand Down
8 changes: 8 additions & 0 deletions ggml/src/ggml-cuda/mmvf.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -961,6 +961,14 @@ bool ggml_cuda_should_use_mmvf(enum ggml_type type, int cc, const int64_t * src0
}
}

// Same one-tile argument as ggml_cuda_should_use_mmvq: below one MMF/MMQ tile the tiled
// kernels launch a single block and leave the device idle, so take the vector kernel
// whatever the type says. src0_ne[1] = 48 is the GDN alpha/beta projection shape.
if (GGML_CUDA_CC_IS_RDNA3_5(cc) && src0_ne[1] < 64 && ne11 <= MMVF_MAX_BATCH_SIZE &&
(type == GGML_TYPE_F32 || type == GGML_TYPE_F16 || type == GGML_TYPE_BF16)) {
return true;
}

switch (type) {
case GGML_TYPE_F32:
if (GGML_CUDA_CC_IS_NVIDIA(cc)) {
Expand Down
Loading