diff --git a/ggml/src/ggml-common.h b/ggml/src/ggml-common.h index 087376502..72a170e40 100644 --- a/ggml/src/ggml-common.h +++ b/ggml/src/ggml-common.h @@ -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 diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index d27d8acb1..7b686bd5e 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -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 struct ggml_cuda_type_traits; @@ -985,6 +1005,13 @@ struct ggml_cuda_type_traits { static constexpr int qi = QI2_0; }; +template<> +struct ggml_cuda_type_traits { + static constexpr int qk = QK_DT3; + static constexpr int qr = QR_DT3; + static constexpr int qi = QI_DT3; +}; + template<> struct ggml_cuda_type_traits { static constexpr int qk = QK4_0; diff --git a/ggml/src/ggml-cuda/convert.cu b/ggml/src/ggml-cuda/convert.cu index 360c614a4..85a7fa71c 100644 --- a/ggml/src/ggml-cuda/convert.cu +++ b/ggml/src/ggml-cuda/convert.cu @@ -461,6 +461,8 @@ to_bf16_cuda_t ggml_get_to_bf16_cuda(ggml_type type) { return dequantize_block_cont_cuda; case GGML_TYPE_Q2_0: return dequantize_block_cont_cuda; + case GGML_TYPE_DT3: + return dequantize_block_cont_cuda; 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; case GGML_TYPE_Q2_0: return dequantize_block_cont_cuda; + case GGML_TYPE_DT3: + return dequantize_block_cont_cuda; 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; case GGML_TYPE_Q2_0: return dequantize_block_cont_cuda; + case GGML_TYPE_DT3: + return dequantize_block_cont_cuda; 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; case GGML_TYPE_Q2_0: return dequantize_block_cuda; + case GGML_TYPE_DT3: + return dequantize_block_cuda; case GGML_TYPE_Q4_0: return dequantize_block_cuda; 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; case GGML_TYPE_Q2_0: return dequantize_block_cuda; + case GGML_TYPE_DT3: + return dequantize_block_cuda; case GGML_TYPE_Q4_0: return dequantize_block_cuda; 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; case GGML_TYPE_Q2_0: return dequantize_block_cuda; + case GGML_TYPE_DT3: + return dequantize_block_cuda; case GGML_TYPE_Q4_0: return dequantize_block_cuda; case GGML_TYPE_Q4_1: diff --git a/ggml/src/ggml-cuda/dequantize.cuh b/ggml/src/ggml-cuda/dequantize.cuh index 651524be7..ae814802e 100644 --- a/ggml/src/ggml-cuda/dequantize.cuh +++ b/ggml/src/ggml-cuda/dequantize.cuh @@ -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; diff --git a/ggml/src/ggml-cuda/getrows.cu b/ggml/src/ggml-cuda/getrows.cu index 6b36b8fbc..264d14725 100644 --- a/ggml/src/ggml-cuda/getrows.cu +++ b/ggml/src/ggml-cuda/getrows.cu @@ -324,6 +324,10 @@ static void ggml_cuda_get_rows_switch_src0_type( get_rows_cuda_q(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(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(src0_d, src1_d, dst_d, ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream); diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index dec619324..1d2713cf5 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -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: