ggml-cuda: widen Q4_K MMVQ weight loads to global_load_b128

ggml/src/ggml-cuda/mmvq.cu

@@ -605,6 +605,15 @@ static __global__ void mul_mat_vec_q(
         // x block quant index when casting the quants to int
         const int kqs = vdr * (tid % (qi/vdr));
 
+#if defined(RDNA3_5)
+        // Approach A weight-hoist: Q4_K, ncols>1, no fusion, single row/block.
+        // Decode the weight group once and reuse it across the ncols activation columns.
+        if constexpr (type == GGML_TYPE_Q4_K && ncols_dst > 1 && !has_fusion && rows_per_cuda_block == 1) {
+            vec_dot_q4_K_q8_1_ncols<ncols_dst>(
+                vx, kbx_offset + kbx, y, kby, stride_col_y, kqs, &tmp[0][0]);
+        } else
+#endif // defined(RDNA3_5)
+        {
 #pragma unroll
         for (int j = 0; j < ncols_dst; ++j) {
 #pragma unroll
@@ -619,6 +628,7 @@ static __global__ void mul_mat_vec_q(
                 }
             }
         }
+        }
     }
 
     __shared__ float tmp_shared[nwarps-1 > 0 ? nwarps-1 : 1][ncols_dst][rows_per_cuda_block][warp_size];

ggml/src/ggml-cuda/vecdotq.cuh

@@ -569,10 +569,14 @@ static __device__ __forceinline__ float vec_dot_q3_K_q8_1_impl_mmq(
     return d3*d8 * sumi;
 }
 
-#define VDR_Q4_K_Q8_1_MMVQ 2
+#if defined(RDNA3_5)
+#define VDR_Q4_K_Q8_1_MMVQ 8   // b128 wide-load path: one thread owns a full 32-byte group
+#else
+#define VDR_Q4_K_Q8_1_MMVQ 2   // upstream default on non-RDNA3.5 targets
+#endif
 #define VDR_Q4_K_Q8_1_MMQ  8
 
-// contiguous v/x values
+// contiguous v/x values (VDR=2 fallback; used on non-RDNA3.5 targets)
 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,
     const uint8_t * __restrict__ m, const half2 & dm4, const float * __restrict__ d8) {
@@ -597,6 +601,98 @@ static __device__ __forceinline__ float vec_dot_q4_K_q8_1_impl_vmmq(
     return dm4f.x*sumf_d - dm4f.y*sumf_m;
 }
 
+// contiguous v/x values
+// 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;
+    float sumf_m = 0.0f;
+
+#pragma unroll
+    for (int i = 0; i < QR4_K; ++i) {
+#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, 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
+        }
+    }
+
+    const float2 dm4f = __half22float2(dm4);
+
+    return dm4f.x*sumf_d - dm4f.y*sumf_m;
+}
+
+#if defined(RDNA3_5)
+// Shared Q4_K b128 (VDR=8) decode helpers, used by both the single-column vec_dot function-
+// pointer target (vec_dot_q4_K_q8_1) and the multi-column weight-hoist (vec_dot_q4_K_q8_1_ncols).
+struct q4_K_weight_vmmq8 {
+    int     vlo[4];
+    int     vhi[4];
+    uint8_t sc[2];
+    uint8_t m[2];
+    half2   dm;
+    int     bq8_offset;
+};
+
+static __device__ __forceinline__ q4_K_weight_vmmq8 load_q4_K_weight_vmmq8(
+    const block_q4_K * __restrict__ bq4_K, const int iqs) {
+
+    q4_K_weight_vmmq8 w;
+    // iqs is in {0,8,16,24}; each thread owns one full 32-byte group.
+    w.bq8_offset = QR4_K * ((iqs/2) / (QI8_1/2)); // {0,2,4,6}; group = bq8_offset/2
+
+    // 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 * w.bq8_offset);
+    *((int4 *) w.vlo) = *((const int4 *) (q4base + 0));
+    *((int4 *) w.vhi) = *((const int4 *) (q4base + 4));
+
+    const uint16_t * scales = (const uint16_t *)bq4_K->scales;
+    uint16_t aux[2];
+    const int j = w.bq8_offset/2;
+    if (j < 2) {
+        aux[0] = scales[j+0] & 0x3f3f;
+        aux[1] = scales[j+2] & 0x3f3f;
+    } else {
+        aux[0] = ((scales[j+2] >> 0) & 0x0f0f) | ((scales[j-2] & 0xc0c0) >> 2);
+        aux[1] = ((scales[j+2] >> 4) & 0x0f0f) | ((scales[j-0] & 0xc0c0) >> 2);
+    }
+    const uint8_t * sc = (const uint8_t *)aux;
+    w.sc[0] = sc[0]; w.sc[1] = sc[1];
+    w.m[0]  = sc[2]; w.m[1]  = sc[3];
+    w.dm = bq4_K->dm;
+    return w;
+}
+
+static __device__ __forceinline__ void load_q4_K_activation_vmmq8(
+    const block_q8_1 * __restrict__ bq8_1, const int bq8_offset,
+    int * __restrict__ ulo, int * __restrict__ uhi, float * __restrict__ d8) {
+
+    // q8_1 blocks are 36 bytes apart (qs 4-byte aligned), so these stay scalar b32.
+#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);
+#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);
+        }
+    }
+}
+
+#endif // defined(RDNA3_5)
+
 // contiguous v/x + u/y values
 static __device__ __forceinline__ float vec_dot_q4_K_q8_1_impl_mmq(
     const int * __restrict__ v, const int * __restrict__ u, const uint8_t * __restrict__ sc,
@@ -1000,18 +1096,24 @@ static __device__ __forceinline__ float vec_dot_q4_K_q8_1(
 
     const block_q4_K * bq4_K = (const block_q4_K *) vbq + kbx;
 
+#if defined(RDNA3_5)
+    // VDR=8: decode the weight group + scales once, then load this column's activation.
+    const q4_K_weight_vmmq8 w = load_q4_K_weight_vmmq8(bq4_K, iqs);
+
+    int   ulo[QR4_K*4];
+    int   uhi[QR4_K*4];
+    float d8[QR4_K];
+    load_q4_K_activation_vmmq8(bq8_1, w.bq8_offset, ulo, uhi, d8);
+
+    return vec_dot_q4_K_q8_1_impl_vmmq8(w.vlo, w.vhi, ulo, uhi, w.sc, w.m, w.dm, d8);
+#else
     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));
 
-    // 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];
@@ -1039,7 +1141,32 @@ static __device__ __forceinline__ float vec_dot_q4_K_q8_1(
     }
 
     return vec_dot_q4_K_q8_1_impl_vmmq(v, u, sc, m, bq4_K->dm, d8);
+#endif // defined(RDNA3_5)
+}
+
+#if defined(RDNA3_5)
+template <int ncols_dst>
+static __device__ __forceinline__ void vec_dot_q4_K_q8_1_ncols(
+    const void * __restrict__ vbq, const int kbx,
+    const block_q8_1 * __restrict__ y, const int kby, const int stride_col_y,
+    const int iqs, float * __restrict__ acc) {
+
+    const block_q4_K * bq4_K = (const block_q4_K *) vbq + kbx;
+    const q4_K_weight_vmmq8 w = load_q4_K_weight_vmmq8(bq4_K, iqs);
+
+#pragma unroll
+    for (int j = 0; j < ncols_dst; ++j) {
+        const block_q8_1 * bq8_1 = &y[j*stride_col_y + kby];
+
+        int   ulo[QR4_K*4];
+        int   uhi[QR4_K*4];
+        float d8[QR4_K];
+        load_q4_K_activation_vmmq8(bq8_1, w.bq8_offset, ulo, uhi, d8);
+
+        acc[j] += vec_dot_q4_K_q8_1_impl_vmmq8(w.vlo, w.vhi, ulo, uhi, w.sc, w.m, w.dm, d8);
+    }
 }
+#endif // defined(RDNA3_5)
 
 static __device__ __forceinline__ float vec_dot_q5_K_q8_1(
     const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {