millaguie/llama.cpp
1
0
Fork 0
Code Issues Pull Requests Actions 32 Packages Projects Releases Wiki Activity

cuda : add DT3 MMVQ kernel

Browse Source 
vec_dot_dt3_q8_1 processes a whole 128-element DT3 block per call
(VDR_DT3_Q8_1_MMVQ = 4, QI_DT3 = 4), i.e. the 4 q8_1 chunks it spans,
with one pair of integer accumulators per chunk:

    sum_j d8[j] * (d1*sumi1[j] + d2*sumi2[j])

The trit decode reuses ggml_cuda_dt3_get_trit with fully unrolled loops,
so all indices and pow3 factors fold into constants; no __byte_perm or
other NVIDIA-only intrinsics. Enables MUL_MAT in supports_op: ncols_dst
<= 8 takes MMVQ, larger falls back to dequantization + cuBLAS (no MMQ
tile kernel yet).
This commit is contained in:
Millaguie
2026-08-10 23:33:14 +02:00
parent 1a1f869f93
commit 0cc5e310c1
3 changed files with 57 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/ggml-cuda.cu
+1
View File
@@ -4908,6 +4908,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
case GGML_TYPE_F16:
case GGML_TYPE_Q1_0:
case GGML_TYPE_Q2_0:
case GGML_TYPE_DT3:
case GGML_TYPE_Q4_0:
case GGML_TYPE_Q4_1:
case GGML_TYPE_Q5_0:
ggml/src/ggml-cuda/mmvq.cu
+8
View File
@@ -11,6 +11,7 @@ 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;
case GGML_TYPE_DT3: return vec_dot_dt3_q8_1;
case GGML_TYPE_Q4_0: return vec_dot_q4_0_q8_1;
case GGML_TYPE_Q4_1: return vec_dot_q4_1_q8_1;
case GGML_TYPE_Q5_0: return vec_dot_q5_0_q8_1;
@@ -40,6 +41,7 @@ static constexpr __host__ __device__ int get_vdr_mmvq(ggml_type type) {
switch (type) {
case GGML_TYPE_Q1_0: return VDR_Q1_0_Q8_1_MMVQ;
case GGML_TYPE_Q2_0: return VDR_Q2_0_Q8_1_MMVQ;
case GGML_TYPE_DT3: return VDR_DT3_Q8_1_MMVQ;
case GGML_TYPE_Q4_0: return VDR_Q4_0_Q8_1_MMVQ;
case GGML_TYPE_Q4_1: return VDR_Q4_1_Q8_1_MMVQ;
case GGML_TYPE_Q5_0: return VDR_Q5_0_Q8_1_MMVQ;
@@ -1018,6 +1020,12 @@ static void mul_mat_vec_q_switch_type(
nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst,
nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream);
break;
case GGML_TYPE_DT3:
mul_mat_vec_q_switch_ncols_dst<GGML_TYPE_DT3>
(vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst,
nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst,
nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream);
break;
case GGML_TYPE_Q4_0:
mul_mat_vec_q_switch_ncols_dst<GGML_TYPE_Q4_0>
(vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst,
ggml/src/ggml-cuda/vecdotq.cuh
+48
View File
@@ -112,6 +112,8 @@ static __device__ __forceinline__ uint32_t unpack_ksigns(const uint8_t v) {
#define VDR_Q2_0_Q8_1_MMVQ 1 // Process one 32-element chunk at a time for parallelism
#define VDR_Q2_0_Q8_1_MMQ 2 // Q2_0 group 64: 128 bits (4 ints) per block, 2 32-element chunks
#define VDR_DT3_Q8_1_MMVQ 4 // DT3: one call processes a whole 128-element block, i.e. all 4 q8_1 chunks it spans
#define VDR_Q4_0_Q8_1_MMVQ 2
#define VDR_Q4_0_Q8_1_MMQ 4
@@ -763,6 +765,52 @@ static __device__ __forceinline__ float vec_dot_q2_0_q8_1(
return d2 * d8 * sumi;
}
static __device__ __forceinline__ float vec_dot_dt3_q8_1(
const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {
const block_dt3 * bq_dt3 = (const block_dt3 *) vbq + kbx;
// DT3: 128 elements as two ternary planes with one scale each, w = d1*t1 + d2*t2.
// One call processes the whole block (VDR_DT3_Q8_1_MMVQ == 4), so iqs is always 0
// and bq8_1 points to the 4 q8_1 blocks the DT3 block spans. All element indices
// below are compile-time constants, so the decode folds into shifts and masks.
GGML_UNUSED(iqs);
int sumi1[4] = {0, 0, 0, 0};
int sumi2[4] = {0, 0, 0, 0};
#pragma unroll
for (int j = 0; j < 4; ++j) {
#pragma unroll
for (int k = 0; k < 8; ++k) {
const int u = get_int_b4(bq8_1[j].qs, k);
int v1 = 0;
int v2 = 0;
#pragma unroll
for (int l = 0; l < 4; ++l) {
const int i = 32*j + 4*k + l;
v1 |= (ggml_cuda_dt3_get_trit(bq_dt3->qs[0], bq_dt3->qh[0], i) & 0xFF) << (8*l);
v2 |= (ggml_cuda_dt3_get_trit(bq_dt3->qs[1], bq_dt3->qh[1], i) & 0xFF) << (8*l);
}
sumi1[j] = ggml_cuda_dp4a(v1, u, sumi1[j]);
sumi2[j] = ggml_cuda_dp4a(v2, u, sumi2[j]);
}
}
const float d1 = bq_dt3->d[0];
const float d2 = bq_dt3->d[1];
float sumf = 0.0f;
#pragma unroll
for (int j = 0; j < 4; ++j) {
const float d8 = __low2float(bq8_1[j].ds);
sumf += d8 * (d1*sumi1[j] + d2*sumi2[j]);
}
return sumf;
}
static __device__ __forceinline__ float vec_dot_q4_0_q8_1(
const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {