wmantly/llama.cpp-xrip
1
0 Fork 0
Code Issues Pull Requests Actions 22 Packages Projects Releases Wiki Activity

ggml-cuda: use DP4A offset for all Q6_K MMVQ paths

Browse Source
This commit is contained in:
xrip committed 2026-09-21 20:46:09 +03:00
1 parent 255a4421fd
commit e49563cdf8
2 files changed
+6 -15
Download Patch File Download Diff File Expand all files Collapse all files

No files matched your search

ggml/src/ggml-cuda/mmvq.cu
+3 -4
View File
@@ -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<true> : vec_dot_q6_K_q8_1<false>;
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;
ggml/src/ggml-cuda/vecdotq.cuh
+3 -11
View File
@@ -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 <bool use_nonsaturating_subtract>
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 <bool use_nonsaturating_subtract>
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<use_nonsaturating_subtract>(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