From 08f7e4fcc589a2a0eb50bfbc5bf42adea5fd477a Mon Sep 17 00:00:00 2001 From: Annie Ren Date: Fri, 7 Aug 2026 12:12:55 -0600 Subject: [PATCH] ggml-cuda: widen Q4_K MMVQ weight loads to global_load_b128 (VDR 2->8) Set VDR_Q4_K_Q8_1_MMVQ to 8 so each thread owns a full 32-byte Q4_K group and loads its weights as two int4 (lowered to global_load_b128), cutting weight-load instructions on the bandwidth-dominant side. Activations stay scalar b32 (block_q8_1 is 36B, not 16B-aligned). Numerically equivalent to the VDR=2 path. Measured +1-2% typical decode on the MMVQ path (gfx1151, GGML_CUDA_DQ_MMV=0), up to +6-7% on some models; Q4_K-only, no effect on the default DQ path. Co-Authored-By: Claude Opus 4 --- ggml/src/ggml-cuda/vecdotq.cuh | 63 +++++++++++++++++++--------------- 1 file changed, 35 insertions(+), 28 deletions(-) diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index 1ac8283437e..d7a803fc9b6 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -569,12 +569,15 @@ static __device__ __forceinline__ float vec_dot_q3_K_q8_1_impl_mmq( return d3*d8 * sumi; } -#define VDR_Q4_K_Q8_1_MMVQ 2 +#define VDR_Q4_K_Q8_1_MMVQ 8 #define VDR_Q4_K_Q8_1_MMQ 8 // contiguous v/x values -static __device__ __forceinline__ float vec_dot_q4_K_q8_1_impl_vmmq( - const int * __restrict__ v, const int * __restrict__ u, const uint8_t * __restrict__ sc, +// VDR=8: one thread owns a full 32-byte Q4_K group (4 slots). vlo/vhi hold the two 16-byte +// weight halves (b128-loaded); ulo/uhi hold the matching q8_1 ints indexed [i*4 + slot]. +static __device__ __forceinline__ float vec_dot_q4_K_q8_1_impl_vmmq8( + const int * __restrict__ vlo, const int * __restrict__ vhi, + const int * __restrict__ ulo, const int * __restrict__ uhi, const uint8_t * __restrict__ sc, const uint8_t * __restrict__ m, const half2 & dm4, const float * __restrict__ d8) { float sumf_d = 0.0f; @@ -582,14 +585,17 @@ static __device__ __forceinline__ float vec_dot_q4_K_q8_1_impl_vmmq( #pragma unroll for (int i = 0; i < QR4_K; ++i) { - const int v0i = (v[0] >> (4*i)) & 0x0F0F0F0F; - const int v1i = (v[1] >> (4*i)) & 0x0F0F0F0F; +#pragma unroll + for (int s = 0; s < 4; ++s) { + const int v0i = (vlo[s] >> (4*i)) & 0x0F0F0F0F; + const int v1i = (vhi[s] >> (4*i)) & 0x0F0F0F0F; - const int dot1 = ggml_cuda_dp4a(v1i, u[2*i+1], ggml_cuda_dp4a(v0i, u[2*i+0], 0)); // SIMD dot product - const int dot2 = ggml_cuda_dp4a(0x01010101, u[2*i+1], ggml_cuda_dp4a(0x01010101, u[2*i+0], 0)); // sum of u + const int dot1 = ggml_cuda_dp4a(v1i, uhi[i*4 + s], ggml_cuda_dp4a(v0i, ulo[i*4 + s], 0)); // SIMD dot product + const int dot2 = ggml_cuda_dp4a(0x01010101, uhi[i*4 + s], ggml_cuda_dp4a(0x01010101, ulo[i*4 + s], 0)); // sum of u - sumf_d += d8[i] * (dot1 * sc[i]); - sumf_m += d8[i] * (dot2 * m[i]); // multiply constant part of q4_K with sum of q8_1 values + sumf_d += d8[i] * (dot1 * sc[i]); + sumf_m += d8[i] * (dot2 * m[i]); // multiply constant part of q4_K with sum of q8_1 values + } } const float2 dm4f = __half22float2(dm4); @@ -1000,21 +1006,16 @@ static __device__ __forceinline__ float vec_dot_q4_K_q8_1( const block_q4_K * bq4_K = (const block_q4_K *) vbq + kbx; - int v[2]; - int u[2*QR4_K]; - float d8[QR4_K]; - - // iqs is in 0,2..30. bq8_offset = iqs/4 -> bq8_offset = 0, 2, 4, 6 - const int bq8_offset = QR4_K * ((iqs/2) / (QI8_1/2)); + // VDR=8: iqs is in {0,8,16,24}; each thread owns one full 32-byte group. + const int bq8_offset = QR4_K * ((iqs/2) / (QI8_1/2)); // {0,2,4,6}; group = bq8_offset/2 - // iqs = 0....3 -> bq8_offset = 0, want q4_offset = 0, 4, 8, 12 - // iqs = 4....7 -> bq8_offset = 2, want q4_offset = 32, 36, 40, 44 - // iqs = 8...11 -> bq8_offset = 4, want q4_offset = 64, 68, 72, 76 - // iqs = 12..15 -> bq8_offset = 6, want q4_offset = 96, 100, 104, 108 - - const int * q4 = (const int *)(bq4_K->qs + 16 * bq8_offset + 4 * ((iqs/2)%4)); - v[0] = q4[0]; - v[1] = q4[4]; + // Weight: load the whole 32-byte group as two 16-byte int4 (-> global_load_b128). + // q4base is 16-byte aligned (qs at +16 in a 144-byte block; 16*bq8_offset is a mult. of 16). + const int * q4base = (const int *)(bq4_K->qs + 16 * bq8_offset); + int vlo[4]; // bytes 0..15 of the group (old v[0] across the 4 slots) + int vhi[4]; // bytes 16..31 of the group (old v[1] across the 4 slots) + *((int4 *) vlo) = *((const int4 *) (q4base + 0)); + *((int4 *) vhi) = *((const int4 *) (q4base + 4)); const uint16_t * scales = (const uint16_t *)bq4_K->scales; uint16_t aux[2]; @@ -1029,16 +1030,22 @@ static __device__ __forceinline__ float vec_dot_q4_K_q8_1( const uint8_t * sc = (const uint8_t *)aux; const uint8_t * m = sc + 2; + // Activations: q8_1 blocks are 36 bytes apart (qs 4-byte aligned), so these stay scalar b32. + int ulo[QR4_K*4]; + int uhi[QR4_K*4]; + float d8[QR4_K]; +#pragma unroll for (int i = 0; i < QR4_K; ++i) { const block_q8_1 * bq8i = bq8_1 + bq8_offset + i; d8[i] = __low2float(bq8i->ds); - - const int * q8 = (const int *)bq8i->qs + ((iqs/2)%4); - u[2*i+0] = q8[0]; - u[2*i+1] = q8[4]; +#pragma unroll + for (int s = 0; s < 4; ++s) { + ulo[i*4 + s] = get_int_b4(bq8i->qs, s); + uhi[i*4 + s] = get_int_b4(bq8i->qs, s + 4); + } } - return vec_dot_q4_K_q8_1_impl_vmmq(v, u, sc, m, bq4_K->dm, d8); + return vec_dot_q4_K_q8_1_impl_vmmq8(vlo, vhi, ulo, uhi, sc, m, bq4_K->dm, d8); } static __device__ __forceinline__ float vec_dot_q5_K_q8_1(