From 0cc5e310c177c3ad5e3ebf50e35aaf5d97ca4f1a Mon Sep 17 00:00:00 2001 From: Millaguie Date: Mon, 10 Aug 2026 15:01:38 +0200 Subject: [PATCH] cuda : add DT3 MMVQ kernel 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). --- ggml/src/ggml-cuda/ggml-cuda.cu | 1 + ggml/src/ggml-cuda/mmvq.cu | 8 ++++++ ggml/src/ggml-cuda/vecdotq.cuh | 48 +++++++++++++++++++++++++++++++++ 3 files changed, 57 insertions(+) diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index f5299ed0f..e7c669185 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -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: diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 0589e65bd..3484ddd5a 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -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 + (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 (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst, diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index 0f039c735..c2c06db2e 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -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) {