From e49563cdf84bc71a9c895ed14d80199631a40fa2 Mon Sep 17 00:00:00 2001 From: xrip Date: Thu, 3 Sep 2026 13:39:06 +0000 Subject: [PATCH] ggml-cuda: use DP4A offset for all Q6_K MMVQ paths --- ggml/src/ggml-cuda/mmvq.cu | 7 +++---- ggml/src/ggml-cuda/vecdotq.cuh | 14 +++----------- 2 files changed, 6 insertions(+), 15 deletions(-) diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 4e7c9617b..8fb790d73 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -37,7 +37,7 @@ static __device__ __forceinline__ void mmvq_prefetch_l2(const void * p) { typedef float (*vec_dot_q_cuda_t)(const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs); -static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type, bool q6_nonsaturating = false) { +static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type) { switch (type) { case GGML_TYPE_Q1_0: return vec_dot_q1_0_q8_1; case GGML_TYPE_Q2_0: return vec_dot_q2_0_q8_1; @@ -53,7 +53,7 @@ static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type, case GGML_TYPE_Q4_K: return vec_dot_q4_K_q8_1; case GGML_TYPE_Q5_K: return vec_dot_q5_K_q8_1; case GGML_TYPE_Q6_K: - return q6_nonsaturating ? vec_dot_q6_K_q8_1 : vec_dot_q6_K_q8_1; + return vec_dot_q6_K_q8_1; case GGML_TYPE_IQ2_XXS: return vec_dot_iq2_xxs_q8_1; case GGML_TYPE_IQ2_XS: return vec_dot_iq2_xs_q8_1; case GGML_TYPE_IQ2_S: return vec_dot_iq2_s_q8_1; @@ -619,8 +619,7 @@ static __global__ void mul_mat_vec_q( constexpr int rows_per_cuda_block = calc_rows_per_block(ncols_dst, table_id, small_k, nwarps); constexpr int warp_size = ggml_cuda_get_physical_warp_size(); - constexpr bool q6_nonsaturating = table_id == MMVQ_PARAMETERS_TURING && ncols_dst == 1; - constexpr vec_dot_q_cuda_t vec_dot_q_cuda = get_vec_dot_q_cuda(type, q6_nonsaturating); + constexpr vec_dot_q_cuda_t vec_dot_q_cuda = get_vec_dot_q_cuda(type); const int tid = warp_size*threadIdx.y + threadIdx.x; const int row0 = rows_per_cuda_block*blockIdx.x; diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index c10dd396c..bd46e3ad1 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -624,7 +624,6 @@ static __device__ __forceinline__ float vec_dot_q5_K_q8_1_impl_mmq( #define VDR_Q6_K_Q8_1_MMQ 8 // contiguous v/x values -template static __device__ __forceinline__ float vec_dot_q6_K_q8_1_impl_mmvq( const int & vl, const int & vh, const int * __restrict__ u, const int8_t * __restrict__ scales, const float & d, const float * __restrict__ d8) { @@ -639,14 +638,8 @@ static __device__ __forceinline__ float vec_dot_q6_K_q8_1_impl_mmvq( const int vih = ((vh >> (4*i)) << 4) & 0x30303030; - int dot; - if constexpr (use_nonsaturating_subtract) { - const int vi = vil | vih; - dot = ggml_cuda_dp4a(vi, u[i], ggml_cuda_dp4a(0xE0E0E0E0, u[i], 0)); - } else { - const int vi = __vsubss4((vil | vih), 0x20202020); // vi = (vil | vih) - 32 - dot = ggml_cuda_dp4a(vi, u[i], 0); - } + const int vi = vil | vih; + const int dot = ggml_cuda_dp4a(vi, u[i], ggml_cuda_dp4a(0xE0E0E0E0, u[i], 0)); sumf += d8[i] * (dot * sc); // SIMD dot product } @@ -1024,7 +1017,6 @@ static __device__ __forceinline__ float vec_dot_q5_K_q8_1( return vec_dot_q5_K_q8_1_impl_vmmq(vl, vh, u, sc, m, bq5_K->dm, d8); } -template static __device__ __forceinline__ float vec_dot_q6_K_q8_1( const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { @@ -1048,7 +1040,7 @@ static __device__ __forceinline__ float vec_dot_q6_K_q8_1( d8[i] = __low2float(bq8_1[bq8_offset + 2*i].ds); } - return vec_dot_q6_K_q8_1_impl_mmvq(vl, vh, u, scales, bq6_K->d, d8); + return vec_dot_q6_K_q8_1_impl_mmvq(vl, vh, u, scales, bq6_K->d, d8); } #define VDR_IQ2_XXS_Q8_1_MMVQ 2