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

cuda : add DT3 dequantization

Browse Source 
Decode one packed ternary plane with the shared ggml_cuda_dt3_get_trit
helper (shifts, masks and a small pow3 table; the uint8_t wrap-around of
the intermediate product is intentional and matches the CPU reference).
The qh bytes hold only 4 trits; their 5th base-3 digit is packing padding
that always decodes to -1 and is never read.

Wires DT3 into the generic dequantize_block templates (to fp32/fp16/bf16,
contiguous and not) and into get_rows, and enables GET_ROWS in
supports_op.
This commit is contained in:
Millaguie
2026-08-10 15:01:28 +02:00
parent 0ecb2a521e
commit 4dfb210207
6 changed files with 60 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-common.h
+3
View File
@@ -99,6 +99,9 @@ typedef sycl::half2 ggml_half2;
#define QI2_0 (QK2_0 / 32)
#define QR2_0 1
#define QI_DT3 (QK_DT3 / 32)
#define QR_DT3 1
#define QI4_0 (QK4_0 / (4 * QR4_0))
#define QR4_0 2
ggml/src/ggml-cuda/common.cuh
+27
View File
@@ -962,6 +962,26 @@ static __device__ __forceinline__ float get_alibi_slope(
return powf(base, exph);
}
// decode element i (0..QK_DT3) of one packed DT3 ternary plane to a trit in {-1, 0, +1}
// layout per plane (see block_dt3): qs[0..16) hold elements m + n*16 (m = 0..16, n = 0..5),
// qs[16..24) hold elements 80 + m + n*8 (m = 0..8, n = 0..5), qh[0..2) hold elements
// 120 + j + n*2 (j = 0..2, n = 0..4) — a qh byte stores only 4 trits, its 5th base-3
// digit is packing padding that always decodes to -1 and must never be read
static __device__ __forceinline__ int ggml_cuda_dt3_get_trit(
const uint8_t * __restrict__ qs, const uint8_t * __restrict__ qh, const int i) {
const uint8_t pow3[5] = {1, 3, 9, 27, 81};
uint8_t q; // the multiplications below wrap around in uint8_t on purpose
if (i < 80) {
q = qs[i % 16] * pow3[i / 16];
} else if (i < 120) {
q = qs[16 + (i - 80) % 8] * pow3[(i - 80) / 8];
} else {
q = qh[(i - 120) % 2] * pow3[(i - 120) / 2];
}
return (int) (((uint16_t) q * 3) >> 8) - 1;
}
template <ggml_type type>
struct ggml_cuda_type_traits;
@@ -985,6 +1005,13 @@ struct ggml_cuda_type_traits<GGML_TYPE_Q2_0> {
static constexpr int qi = QI2_0;
};
template<>
struct ggml_cuda_type_traits<GGML_TYPE_DT3> {
static constexpr int qk = QK_DT3;
static constexpr int qr = QR_DT3;
static constexpr int qi = QI_DT3;
};
template<>
struct ggml_cuda_type_traits<GGML_TYPE_Q4_0> {
static constexpr int qk = QK4_0;
ggml/src/ggml-cuda/convert.cu
+12
View File
@@ -461,6 +461,8 @@ to_bf16_cuda_t ggml_get_to_bf16_cuda(ggml_type type) {
return dequantize_block_cont_cuda<QK1_0, QR1_0, dequantize_q1_0>;
case GGML_TYPE_Q2_0:
return dequantize_block_cont_cuda<QK2_0, QR2_0, dequantize_q2_0>;
case GGML_TYPE_DT3:
return dequantize_block_cont_cuda<QK_DT3, QR_DT3, dequantize_dt3>;
case GGML_TYPE_Q4_0:
return dequantize_row_q4_0_cuda;
case GGML_TYPE_Q4_1:
@@ -518,6 +520,8 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) {
return dequantize_block_cont_cuda<QK1_0, QR1_0, dequantize_q1_0>;
case GGML_TYPE_Q2_0:
return dequantize_block_cont_cuda<QK2_0, QR2_0, dequantize_q2_0>;
case GGML_TYPE_DT3:
return dequantize_block_cont_cuda<QK_DT3, QR_DT3, dequantize_dt3>;
case GGML_TYPE_Q4_0:
return dequantize_row_q4_0_cuda;
case GGML_TYPE_Q4_1:
@@ -578,6 +582,8 @@ to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) {
return dequantize_block_cont_cuda<QK1_0, QR1_0, dequantize_q1_0>;
case GGML_TYPE_Q2_0:
return dequantize_block_cont_cuda<QK2_0, QR2_0, dequantize_q2_0>;
case GGML_TYPE_DT3:
return dequantize_block_cont_cuda<QK_DT3, QR_DT3, dequantize_dt3>;
case GGML_TYPE_Q4_0:
return dequantize_row_q4_0_cuda;
case GGML_TYPE_Q4_1:
@@ -637,6 +643,8 @@ to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type) {
return dequantize_block_cuda<QK1_0, QR1_0, dequantize_q1_0>;
case GGML_TYPE_Q2_0:
return dequantize_block_cuda<QK2_0, QR2_0, dequantize_q2_0>;
case GGML_TYPE_DT3:
return dequantize_block_cuda<QK_DT3, QR_DT3, dequantize_dt3>;
case GGML_TYPE_Q4_0:
return dequantize_block_cuda<QK4_0, QR4_0, dequantize_q4_0>;
case GGML_TYPE_Q4_1:
@@ -662,6 +670,8 @@ to_bf16_nc_cuda_t ggml_get_to_bf16_nc_cuda(ggml_type type) {
return dequantize_block_cuda<QK1_0, QR1_0, dequantize_q1_0>;
case GGML_TYPE_Q2_0:
return dequantize_block_cuda<QK2_0, QR2_0, dequantize_q2_0>;
case GGML_TYPE_DT3:
return dequantize_block_cuda<QK_DT3, QR_DT3, dequantize_dt3>;
case GGML_TYPE_Q4_0:
return dequantize_block_cuda<QK4_0, QR4_0, dequantize_q4_0>;
case GGML_TYPE_Q4_1:
@@ -687,6 +697,8 @@ to_fp32_nc_cuda_t ggml_get_to_fp32_nc_cuda(ggml_type type) {
return dequantize_block_cuda<QK1_0, QR1_0, dequantize_q1_0>;
case GGML_TYPE_Q2_0:
return dequantize_block_cuda<QK2_0, QR2_0, dequantize_q2_0>;
case GGML_TYPE_DT3:
return dequantize_block_cuda<QK_DT3, QR_DT3, dequantize_dt3>;
case GGML_TYPE_Q4_0:
return dequantize_block_cuda<QK4_0, QR4_0, dequantize_q4_0>;
case GGML_TYPE_Q4_1:
ggml/src/ggml-cuda/dequantize.cuh
+13
View File
@@ -43,6 +43,19 @@ static __device__ __forceinline__ void dequantize_q2_0(const void * vx, const in
v.y = (c1 - 1) * d;
}
static __device__ __forceinline__ void dequantize_dt3(const void * vx, const int64_t ib, const int iqs, float2 & v){
const block_dt3 * x = (const block_dt3 *) vx;
// DT3: two packed ternary planes with one scale each, w = d1*t1 + d2*t2
const float d1 = x[ib].d[0];
const float d2 = x[ib].d[1];
v.x = d1*ggml_cuda_dt3_get_trit(x[ib].qs[0], x[ib].qh[0], iqs + 0) +
d2*ggml_cuda_dt3_get_trit(x[ib].qs[1], x[ib].qh[1], iqs + 0);
v.y = d1*ggml_cuda_dt3_get_trit(x[ib].qs[0], x[ib].qh[0], iqs + 1) +
d2*ggml_cuda_dt3_get_trit(x[ib].qs[1], x[ib].qh[1], iqs + 1);
}
static __device__ __forceinline__ void dequantize_q4_0(const void * vx, const int64_t ib, const int iqs, float2 & v){
const block_q4_0 * x = (const block_q4_0 *) vx;
ggml/src/ggml-cuda/getrows.cu
+4
View File
@@ -324,6 +324,10 @@ static void ggml_cuda_get_rows_switch_src0_type(
get_rows_cuda_q<QK2_0, QR2_0, dequantize_q2_0>(src0_d, src1_d, dst_d,
ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);
break;
case GGML_TYPE_DT3:
get_rows_cuda_q<QK_DT3, QR_DT3, dequantize_dt3>(src0_d, src1_d, dst_d,
ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);
break;
case GGML_TYPE_Q4_0:
get_rows_cuda_q<QK4_0, QR4_0, dequantize_q4_0>(src0_d, src1_d, dst_d,
ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);
ggml/src/ggml-cuda/ggml-cuda.cu
+1
View File
@@ -4947,6 +4947,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
case GGML_TYPE_I32:
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: