From 6f65493dde404ec7054a2a7c3d455e4e155a3d26 Mon Sep 17 00:00:00 2001 From: xrip Date: Mon, 7 Sep 2026 19:21:09 +0300 Subject: [PATCH] cuda: drop SM75 Q4_K decode reuse, superseded by upstream branchless unpack Same story as the packed scale loads. The Q4_K reuse path decoded a block once and applied it across ncols_dst columns, which beat the old branchy generic path because that path re-ran the scale unpack per column. Upstream's branchless rewrite hoists the unpack out of the loop, and the generic path is now the faster of the two. Measured on a CMP 50HX (cc 7.5), test-backend-ops perf, medians of 5 reps after a discarded warmup, two independent runs: reuse disabled vs enabled, q4_K n=2..4: -0.52% and -0.54% (faster without) n=5..8, identical code in both builds: -0.02% (noise check) widening the gate to n<=8 instead: +0.80% (worse) 14 of 15 cases at n=2..4 favour removal, so the effect is consistent rather than drift. The IQ4_XS reuse path is deliberately kept: upstream's change only touched Q4_K/Q5_K scale unpacking, IQ4_XS decodes via kvalues_iq4nl table lookups and its generic path did not improve. Measured with GGML_SM75_IQ4_REUSE=0, removing it costs +1.14% at n=2..4, so it stays. --- ggml/src/ggml-cuda/mmvq.cu | 24 --------------- ggml/src/ggml-cuda/vecdotq.cuh | 56 ---------------------------------- 2 files changed, 80 deletions(-) diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 824cb8bc9..7fba66a3d 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -745,30 +745,6 @@ static __global__ void mul_mat_vec_q( } #endif -#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_TURING - if constexpr (type == GGML_TYPE_Q4_K && ncols_dst >= 2 && ncols_dst <= 4) { -#pragma unroll - for (int i = 0; i < rows_per_cuda_block; ++i) { - const int bx = kbx_offset + i * stride_row_x + kbx; - const sm75_q4_K_decoded w = sm75_q4_K_decode(vx, bx, kqs); -#pragma unroll - for (int j = 0; j < ncols_dst; ++j) { - tmp[j][i] += sm75_q4_K_dot(w, &y[j * stride_col_y + kby], kqs); - } - if constexpr (has_fusion) { - if (use_gate) { - const sm75_q4_K_decoded gate = sm75_q4_K_decode(vgate, bx, kqs); -#pragma unroll - for (int j = 0; j < ncols_dst; ++j) { - tmp_gate[j][i] += sm75_q4_K_dot(gate, &y[j * stride_col_y + kby], kqs); - } - } - } - } - continue; - } -#endif - #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_DGX_SPARK // start the next iterations' weight loads early if constexpr (mmvq_should_prefetch(type)) { diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index 42c9ee014..bd46e3ad1 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -916,62 +916,6 @@ static __device__ __forceinline__ float vec_dot_q3_K_q8_1( return vec_dot_q3_K_q8_1_impl_mmvq(vl, vh, u, bq3_K->scales, scale_offset, d, d8); } -#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ == GGML_CUDA_CC_TURING -struct sm75_q4_K_decoded { - int values[4]; - uint16_t scales; - uint16_t mins; - float2 dm; -}; - -static __device__ __forceinline__ sm75_q4_K_decoded sm75_q4_K_decode(const void * vx, int kbx, int iqs) { - static_assert(QR4_K == 2 && QI8_1 == 8 && VDR_Q4_K_Q8_1_MMVQ == 2, "Q4_K reuse expects two words per lane"); - const block_q4_K * b = (const block_q4_K *) vx + kbx; - const int offset = QR4_K * ((iqs / 2) / (QI8_1 / 2)); - const int * q = (const int *) (b->qs + 16 * offset + 4 * ((iqs / 2) % 4)); - sm75_q4_K_decoded w; -#pragma unroll - for (int i = 0; i < QR4_K; ++i) { - w.values[2 * i] = (q[0] >> (4 * i)) & 0x0f0f0f0f; - w.values[2 * i + 1] = (q[4] >> (4 * i)) & 0x0f0f0f0f; - } - const uint32_t * scales = (const uint32_t *) b->scales; - const int j = offset / 2; - const int shift = 16 * (j & 1); - const uint16_t s0 = scales[0] >> shift; - const uint16_t s1 = scales[1] >> shift; - if (j < 2) { - w.scales = s0 & 0x3f3f; - w.mins = s1 & 0x3f3f; - } else { - const uint16_t s2 = scales[2] >> shift; - w.scales = (s2 & 0x0f0f) | ((s0 & 0xc0c0) >> 2); - w.mins = ((s2 >> 4) & 0x0f0f) | ((s1 & 0xc0c0) >> 2); - } - w.dm = __half22float2(b->dm); - return w; -} - -static __device__ __forceinline__ float sm75_q4_K_dot(const sm75_q4_K_decoded & w, const block_q8_1 * y, int iqs) { - const int offset = QR4_K * ((iqs / 2) / (QI8_1 / 2)); - float sum_d = 0.0f; - float sum_m = 0.0f; -#pragma unroll - for (int i = 0; i < QR4_K; ++i) { - const block_q8_1 * b = y + offset + i; - const int * q = (const int *) b->qs + ((iqs / 2) % 4); - const int dot = ggml_cuda_dp4a(w.values[2 * i + 1], q[4], ggml_cuda_dp4a(w.values[2 * i], q[0], 0)); - const int sum = ggml_cuda_dp4a(0x01010101, q[4], ggml_cuda_dp4a(0x01010101, q[0], 0)); - const int scale = (w.scales >> (8 * i)) & 0xff; - const int min = (w.mins >> (8 * i)) & 0xff; - const float d = __low2float(b->ds); - sum_d += d * (dot * scale); - sum_m += d * (sum * min); - } - return w.dm.x * sum_d - w.dm.y * sum_m; -} -#endif - static __device__ __forceinline__ float vec_dot_q4_K_q8_1( const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {