diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h index 5c04f2c81874..8fa52d812754 100644 --- a/ggml/include/ggml.h +++ b/ggml/include/ggml.h @@ -433,7 +433,8 @@ extern "C" { // Prism-private Q2_0 at group size 128 (upstream Q2_0 is group 64). High id so it // slots above upstream types; type_traits is sized to COUNT (143) with 43..141 unused. GGML_TYPE_PQ2_0 = 142, - GGML_TYPE_COUNT = 143, + GGML_TYPE_PTQ1_0 = 143, // Prism-private ternary, group 128 + GGML_TYPE_COUNT = 144, }; // precision @@ -479,6 +480,7 @@ extern "C" { GGML_FTYPE_MOSTLY_Q1_0 = 27, // except 1d tensors GGML_FTYPE_MOSTLY_Q2_0 = 28, // except 1d tensors GGML_FTYPE_MOSTLY_PQ2_0 = 128, // except 1d tensors (Prism-private group-128 Q2_0) + GGML_FTYPE_MOSTLY_PTQ1_0 = 129, // except 1d tensors (Prism-private group-128 ternary) }; // available tensor operations: diff --git a/ggml/src/ggml-common.h b/ggml/src/ggml-common.h index 7cde4dd68bf2..65075cad860c 100644 --- a/ggml/src/ggml-common.h +++ b/ggml/src/ggml-common.h @@ -101,6 +101,8 @@ typedef sycl::half2 ggml_half2; #define QI_PQ2_0 (QK_PQ2_0 / 32) #define QR_PQ2_0 1 +#define QI_PTQ1_0 (QK_PTQ1_0 / 32) +#define QR_PTQ1_0 1 #define QI4_0 (QK4_0 / (4 * QR4_0)) @@ -204,6 +206,19 @@ typedef struct { } block_pq2_0; static_assert(sizeof(block_pq2_0) == sizeof(ggml_half) + QK_PQ2_0 / 4, "wrong pq2_0 block size/padding"); +// PTQ1_0: Prism-private ternary at group size 128. Same base-3 trit packing as +// upstream TQ1_0 (type 34) but one fp16 scale per 128 weights instead of per 256. +// 1.75 bpw vs PQ2_0's 2.125, and lossless for checkpoints that are already ternary +// at group 128 -- TQ1_0 cannot represent those, because a 256-wide scale has to +// discard one of the two group scales it straddles. +#define QK_PTQ1_0 128 +typedef struct { + uint8_t qs[(QK_PTQ1_0 - 4*QK_PTQ1_0/64)/5]; // 24 B, 5 trits per byte -> 120 values + uint8_t qh[QK_PTQ1_0/64]; // 2 B, 4 trits per byte -> 8 values + ggml_half d; // scale +} block_ptq1_0; +static_assert(sizeof(block_ptq1_0) == sizeof(ggml_half) + QK_PTQ1_0/64 + (QK_PTQ1_0 - 4*QK_PTQ1_0/64)/5, "wrong ptq1_0 block size/padding"); + #define QK4_0 32 typedef struct { ggml_half d; // delta diff --git a/ggml/src/ggml-cpu/arch-fallback.h b/ggml/src/ggml-cpu/arch-fallback.h index f9f4930fa452..d283c2ed51be 100644 --- a/ggml/src/ggml-cpu/arch-fallback.h +++ b/ggml/src/ggml-cpu/arch-fallback.h @@ -18,6 +18,7 @@ #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0 #define ggml_vec_dot_q1_0_q8_0_generic ggml_vec_dot_q1_0_q8_0 #define ggml_vec_dot_pq2_0_q8_0_generic ggml_vec_dot_pq2_0_q8_0 +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 #define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0 #define ggml_vec_dot_tq1_0_q8_K_generic ggml_vec_dot_tq1_0_q8_K #define ggml_vec_dot_tq2_0_q8_K_generic ggml_vec_dot_tq2_0_q8_K @@ -79,6 +80,8 @@ #define ggml_gemm_q1_0_4x8_q8_0_generic ggml_gemm_q1_0_4x8_q8_0 #define ggml_gemm_pq2_0_4x8_q8_0_generic ggml_gemm_pq2_0_4x8_q8_0 #elif defined(__aarch64__) || defined(__arm__) || defined(_M_ARM) || defined(_M_ARM64) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 // repack.cpp #define ggml_quantize_mat_q8_K_4x4_generic ggml_quantize_mat_q8_K_4x4 #define ggml_quantize_mat_q8_K_4x8_generic ggml_quantize_mat_q8_K_4x8 @@ -91,6 +94,8 @@ #define ggml_gemm_q2_K_8x8_q8_K_generic ggml_gemm_q2_K_8x8_q8_K #define ggml_gemm_pq2_0_4x8_q8_0_generic ggml_gemm_pq2_0_4x8_q8_0 #elif defined(__x86_64__) || defined(__i386__) || defined(_M_IX86) || defined(_M_X64) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 // quants.c #define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0 // repack.cpp @@ -121,6 +126,8 @@ #define ggml_gemm_q8_0_4x8_q8_0_generic ggml_gemm_q8_0_4x8_q8_0 #define ggml_gemm_q1_0_4x4_q8_0_generic ggml_gemm_q1_0_4x4_q8_0 #elif defined(__POWERPC__) || defined(__powerpc__) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 // ref: https://github.com/ggml-org/llama.cpp/pull/14146#issuecomment-2972561679 // quants.c #define quantize_row_q8_K_generic quantize_row_q8_K @@ -175,6 +182,8 @@ #define ggml_gemm_q1_0_4x8_q8_0_generic ggml_gemm_q1_0_4x8_q8_0 #define ggml_gemm_pq2_0_4x8_q8_0_generic ggml_gemm_pq2_0_4x8_q8_0 #elif defined(__loongarch64) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 // quants.c #define quantize_row_q8_K_generic quantize_row_q8_K #define ggml_vec_dot_tq1_0_q8_K_generic ggml_vec_dot_tq1_0_q8_K @@ -229,6 +238,8 @@ #define ggml_gemm_q1_0_4x8_q8_0_generic ggml_gemm_q1_0_4x8_q8_0 #define ggml_gemm_pq2_0_4x8_q8_0_generic ggml_gemm_pq2_0_4x8_q8_0 #elif defined(__riscv) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 // quants.c #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0 #define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0 @@ -276,6 +287,8 @@ #define ggml_gemm_q1_0_4x8_q8_0_generic ggml_gemm_q1_0_4x8_q8_0 #define ggml_gemm_pq2_0_4x8_q8_0_generic ggml_gemm_pq2_0_4x8_q8_0 #elif defined(__s390x__) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 // quants.c #define quantize_row_q8_K_generic quantize_row_q8_K #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0 @@ -336,6 +349,8 @@ #define ggml_gemm_q1_0_4x8_q8_0_generic ggml_gemm_q1_0_4x8_q8_0 #define ggml_gemm_pq2_0_4x8_q8_0_generic ggml_gemm_pq2_0_4x8_q8_0 #elif defined(__wasm__) +// PTQ1_0 currently has only the generic vec_dot; alias it here until a SIMD version lands +#define ggml_vec_dot_ptq1_0_q8_0_generic ggml_vec_dot_ptq1_0_q8_0 // quants.c #define ggml_vec_dot_tq1_0_q8_K_generic ggml_vec_dot_tq1_0_q8_K #define ggml_vec_dot_tq2_0_q8_K_generic ggml_vec_dot_tq2_0_q8_K diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c index 79c49c5e4d3b..7a1bcdf355b3 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -242,6 +242,12 @@ static const struct ggml_type_traits_cpu type_traits_cpu[GGML_TYPE_COUNT] = { .vec_dot_type = GGML_TYPE_Q8_0, .nrows = 1, }, + [GGML_TYPE_PTQ1_0] = { + .from_float = quantize_row_ptq1_0, + .vec_dot = ggml_vec_dot_ptq1_0_q8_0, + .vec_dot_type = GGML_TYPE_Q8_0, + .nrows = 1, + }, [GGML_TYPE_Q4_0] = { .from_float = quantize_row_q4_0, .vec_dot = ggml_vec_dot_q4_0_q8_0, diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index f5bcbfb6389e..35697d0787d7 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -667,6 +667,7 @@ void ggml_compute_forward_add( case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -1119,6 +1120,7 @@ void ggml_compute_forward_add1( case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -1251,6 +1253,7 @@ void ggml_compute_forward_acc( case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -4522,6 +4525,7 @@ void ggml_compute_forward_out_prod( case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -4799,6 +4803,7 @@ void ggml_compute_forward_set( case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -5025,6 +5030,7 @@ void ggml_compute_forward_get_rows( case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -5783,6 +5789,7 @@ void ggml_compute_forward_clamp( case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: diff --git a/ggml/src/ggml-cpu/quants.c b/ggml/src/ggml-cpu/quants.c index 1971d57a0224..3e275dc70bf3 100644 --- a/ggml/src/ggml-cpu/quants.c +++ b/ggml/src/ggml-cpu/quants.c @@ -34,6 +34,10 @@ void quantize_row_pq2_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, i quantize_row_pq2_0_ref(x, y, k); } +void quantize_row_ptq1_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { + quantize_row_ptq1_0_ref(x, y, k); +} + void quantize_row_q4_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { quantize_row_q4_0_ref(x, y, k); } @@ -274,6 +278,74 @@ void ggml_vec_dot_pq2_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, *s = sumf; } +// PTQ1_0 x Q8_0. The trits are stored base-3 interleaved rather than in element +// order, so decode a block into element order first using the same traversal as +// dequantize_row_ptq1_0 -- that keeps the two provably in step. Four Q8_0 blocks +// cover one 128-wide PTQ1_0 block. +void ggml_vec_dot_ptq1_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { + const int qk = QK_PTQ1_0; + const int nb = n / qk; + + assert(n % qk == 0); + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + + const block_ptq1_0 * GGML_RESTRICT x = vx; + const block_q8_0 * GGML_RESTRICT y = vy; + + static const uint8_t pow3[6] = {1, 3, 9, 27, 81, 243}; + static const size_t stages[3] = {32, 16, 8}; + + float sumf = 0.0f; + + for (int i = 0; i < nb; i++) { + int8_t q[QK_PTQ1_0]; + int o = 0; + + size_t j = 0; + for (size_t st = 0; st < 3; ++st) { + const size_t c = stages[st]; + for (; j + c <= sizeof(x->qs); j += c) { + for (size_t nn = 0; nn < 5; ++nn) { + for (size_t m = 0; m < c; ++m) { + const uint8_t v = x[i].qs[j + m] * pow3[nn]; + const int16_t xi = ((uint16_t) v * 3) >> 8; + q[o++] = (int8_t) (xi - 1); + } + } + } + } + for (size_t nn = 0; nn < 4; ++nn) { + for (size_t h = 0; h < sizeof(x->qh); ++h) { + const uint8_t v = x[i].qh[h] * pow3[nn]; + const int16_t xi = ((uint16_t) v * 3) >> 8; + q[o++] = (int8_t) (xi - 1); + } + } + assert(o == QK_PTQ1_0); + + const float d0 = GGML_CPU_FP16_TO_FP32(x[i].d); + float sumi = 0.0f; + + for (int k = 0; k < 4; k++) { + const block_q8_0 * GGML_RESTRICT yb = &y[i * 4 + k]; + const float d1 = GGML_CPU_FP16_TO_FP32(yb->d); + int sumi_block = 0; + for (int b = 0; b < 32; ++b) { + sumi_block += (int) q[k*32 + b] * (int) yb->qs[b]; + } + sumi += d1 * sumi_block; + } + + sumf += d0 * sumi; + } + + *s = sumf; +} + void ggml_vec_dot_q4_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { const int qk = QK8_0; const int nb = n / qk; diff --git a/ggml/src/ggml-cpu/quants.h b/ggml/src/ggml-cpu/quants.h index 2cac90c6b0a0..f00f35c3fe85 100644 --- a/ggml/src/ggml-cpu/quants.h +++ b/ggml/src/ggml-cpu/quants.h @@ -15,6 +15,7 @@ extern "C" { void quantize_row_q1_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q2_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_pq2_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_ptq1_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q4_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q4_1(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q5_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); @@ -42,6 +43,7 @@ void quantize_row_iq4_xs (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, void ggml_vec_dot_q1_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_q2_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_pq2_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); +void ggml_vec_dot_ptq1_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_q4_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_q4_1_q8_1(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_q5_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); @@ -76,6 +78,7 @@ void quantize_row_q8_1_generic(const float * GGML_RESTRICT x, void * GGML_RESTRI void quantize_row_q8_K_generic(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void ggml_vec_dot_q1_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_pq2_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); +void ggml_vec_dot_ptq1_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_q2_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_q4_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); void ggml_vec_dot_q4_1_q8_1_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc); diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index b2d82beaf7c8..ef929d3d7842 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -992,6 +992,40 @@ struct ggml_cuda_type_traits { static constexpr int qi = QI_PQ2_0; }; +// PTQ1_0 packs trits base-3, five per byte, so element order is not positional. It +// follows the CPU codec exactly: a 16-byte chunk of qs, then an 8-byte chunk, then qh +// at four trits per byte. Trits come out by the base-3 remainder recurrence +// t = (v*3)>>8 with v = (v*3)&0xFF, two integer ops per step and no table. +static __device__ __forceinline__ int ptq1_0_trit(const block_ptq1_0 * x, const int e) { + uint8_t b; + int n; + if (e < 80) { // qs[0..15], chunk of 16 + b = x->qs[e & 15]; n = e >> 4; + } else if (e < 120) { // qs[16..23], chunk of 8 + const int t = e - 80; + b = x->qs[16 + (t & 7)]; n = t >> 3; + } else { // qh[0..1], four trits per byte + const int t = e - 120; + b = x->qh[t & 1]; n = t >> 1; + } + + uint32_t v = b; +#pragma unroll + for (int i = 0; i < 4; ++i) { + if (i < n) { + v = (v * 3) & 0xFF; + } + } + return (int) ((v * 3) >> 8) - 1; +} + +template<> +struct ggml_cuda_type_traits { + static constexpr int qk = QK_PTQ1_0; + static constexpr int qr = QR_PTQ1_0; + static constexpr int qi = QI_PTQ1_0; +}; + 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 807f2caa73ff..f271c00912e9 100644 --- a/ggml/src/ggml-cuda/convert.cu +++ b/ggml/src/ggml-cuda/convert.cu @@ -463,6 +463,8 @@ to_bf16_cuda_t ggml_get_to_bf16_cuda(ggml_type type) { return dequantize_block_cont_cuda; case GGML_TYPE_PQ2_0: return dequantize_block_cont_cuda; + case GGML_TYPE_PTQ1_0: + return dequantize_block_cont_cuda; case GGML_TYPE_Q4_0: return dequantize_row_q4_0_cuda; case GGML_TYPE_Q4_1: @@ -522,6 +524,8 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) { return dequantize_block_cont_cuda; case GGML_TYPE_PQ2_0: return dequantize_block_cont_cuda; + case GGML_TYPE_PTQ1_0: + return dequantize_block_cont_cuda; case GGML_TYPE_Q4_0: return dequantize_row_q4_0_cuda; case GGML_TYPE_Q4_1: @@ -584,6 +588,8 @@ to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) { return dequantize_block_cont_cuda; case GGML_TYPE_PQ2_0: return dequantize_block_cont_cuda; + case GGML_TYPE_PTQ1_0: + return dequantize_block_cont_cuda; case GGML_TYPE_Q4_0: return dequantize_row_q4_0_cuda; case GGML_TYPE_Q4_1: @@ -645,6 +651,8 @@ to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type) { return dequantize_block_cuda; case GGML_TYPE_PQ2_0: return dequantize_block_cuda; + case GGML_TYPE_PTQ1_0: + return dequantize_block_cuda; case GGML_TYPE_Q4_0: return dequantize_block_cuda; case GGML_TYPE_Q4_1: @@ -672,6 +680,8 @@ to_bf16_nc_cuda_t ggml_get_to_bf16_nc_cuda(ggml_type type) { return dequantize_block_cuda; case GGML_TYPE_PQ2_0: return dequantize_block_cuda; + case GGML_TYPE_PTQ1_0: + return dequantize_block_cuda; case GGML_TYPE_Q4_0: return dequantize_block_cuda; case GGML_TYPE_Q4_1: @@ -699,6 +709,8 @@ to_fp32_nc_cuda_t ggml_get_to_fp32_nc_cuda(ggml_type type) { return dequantize_block_cuda; case GGML_TYPE_PQ2_0: return dequantize_block_cuda; + case GGML_TYPE_PTQ1_0: + 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 e3f130f1c8ab..7a1deec8b829 100644 --- a/ggml/src/ggml-cuda/dequantize.cuh +++ b/ggml/src/ggml-cuda/dequantize.cuh @@ -43,6 +43,16 @@ static __device__ __forceinline__ void dequantize_q2_0(const void * vx, const in v.y = (c1 - 1) * d; } + +static __device__ __forceinline__ void dequantize_ptq1_0(const void * vx, const int64_t ib, const int iqs, float2 & v){ + const block_ptq1_0 * x = (const block_ptq1_0 *) vx; + + const float d = x[ib].d; + + v.x = ptq1_0_trit(&x[ib], iqs + 0) * d; + v.y = ptq1_0_trit(&x[ib], iqs + 1) * d; +} + static __device__ __forceinline__ void dequantize_pq2_0(const void * vx, const int64_t ib, const int iqs, float2 & v){ const block_pq2_0 * x = (const block_pq2_0 *) vx; diff --git a/ggml/src/ggml-cuda/getrows.cu b/ggml/src/ggml-cuda/getrows.cu index 49f5bd28dfac..6b4453bc69ce 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_PTQ1_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); + break; case GGML_TYPE_Q2_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 9178d32fe717..dedf60b30d80 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -4988,6 +4988,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -5028,6 +5029,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: 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 768b2394abc3..ede8a65e0154 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -13,6 +13,7 @@ static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type 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_PQ2_0: return vec_dot_pq2_0_q8_1; + case GGML_TYPE_PTQ1_0: return vec_dot_ptq1_0_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; @@ -43,6 +44,7 @@ static constexpr __host__ __device__ int get_vdr_mmvq(ggml_type 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_PQ2_0: return VDR_PQ2_0_Q8_1_MMVQ; + case GGML_TYPE_PTQ1_0: return VDR_PTQ1_0_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; @@ -663,24 +665,52 @@ static __global__ void mul_mat_vec_q( // x block quant index when casting the quants to int const int kqs = vdr * (tid % (qi/vdr)); -#pragma unroll - for (int j = 0; j < ncols_dst; ++j) { -#pragma unroll +#if !defined(GGML_USE_HIP) + if constexpr (type == GGML_TYPE_PTQ1_0 && ncols_dst > 1 && ncols_dst <= 3) { +# pragma unroll for (int i = 0; i < rows_per_cuda_block; ++i) { - tmp[j][i] += vec_dot_q_cuda( - vx, &y[j*stride_col_y + kby], kbx_offset + i*stride_row_x + kbx, kqs); + float dots[ncols_dst]; + vec_dot_ptq1_0_q8_1_multi(vx, &y[kby], kbx_offset + i * stride_row_x + kbx, kqs, + stride_col_y, dots); +# pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + tmp[j][i] += dots[j]; + } + if constexpr (has_fusion) { if (use_gate) { - tmp_gate[j][i] += vec_dot_q_cuda( - vgate, &y[j*stride_col_y + kby], kbx_offset + i*stride_row_x + kbx, kqs); + vec_dot_ptq1_0_q8_1_multi(vgate, &y[kby], kbx_offset + i * stride_row_x + kbx, kqs, + stride_col_y, dots); +# pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + tmp_gate[j][i] += dots[j]; + } + } + } + } + } else +#endif + { +#pragma unroll + for (int j = 0; j < ncols_dst; ++j) { +#pragma unroll + for (int i = 0; i < rows_per_cuda_block; ++i) { + tmp[j][i] += + vec_dot_q_cuda(vx, &y[j * stride_col_y + kby], kbx_offset + i * stride_row_x + kbx, kqs); + if constexpr (has_fusion) { + if (use_gate) { + tmp_gate[j][i] += vec_dot_q_cuda(vgate, &y[j * stride_col_y + kby], + kbx_offset + i * stride_row_x + kbx, kqs); + } } } } } } - __shared__ float tmp_shared[nwarps-1 > 0 ? nwarps-1 : 1][ncols_dst][rows_per_cuda_block][warp_size]; - [[maybe_unused]] __shared__ float tmp_shared_gate[(has_fusion && (nwarps-1 > 0)) ? nwarps-1 : 1][ncols_dst][rows_per_cuda_block][warp_size]; + __shared__ float tmp_shared[nwarps - 1 > 0 ? nwarps - 1 : 1][ncols_dst][rows_per_cuda_block][warp_size]; + [[maybe_unused]] __shared__ float tmp_shared_gate[(has_fusion && (nwarps - 1 > 0)) ? nwarps - 1 : 1][ncols_dst] + [rows_per_cuda_block][warp_size]; if (threadIdx.y > 0) { #pragma unroll @@ -1120,6 +1150,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_PTQ1_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, + 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_Q2_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 f6b1acd591c2..fa6cdabd5d84 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -113,6 +113,7 @@ static __device__ __forceinline__ uint32_t unpack_ksigns(const uint8_t v) { #define VDR_Q2_0_Q8_1_MMQ 2 // Q2_0 group 64: 128 bits (4 ints) per block, 2 32-element chunks #define VDR_PQ2_0_Q8_1_MMVQ 1 // one 32-element chunk at a time (same per-chunk codec as Q2_0) +#define VDR_PTQ1_0_Q8_1_MMVQ 4 // whole 128 block per call: keeps the byte walk uniform across lanes #define VDR_PQ2_0_Q8_1_MMQ 2 // Q2_0 group 128: 4 32-element chunks per block #define VDR_Q4_0_Q8_1_MMVQ 2 @@ -802,9 +803,158 @@ static __device__ __forceinline__ float vec_dot_q2_0_q8_1( return d2 * d8 * sumi; } -static __device__ __forceinline__ float vec_dot_pq2_0_q8_1( - const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { +#if !defined(GGML_USE_HIP) +template +static __device__ __forceinline__ void vec_dot_ptq1_0_q8_1_multi(const void * __restrict__ vbq, + const block_q8_1 * __restrict__ bq8_1, + const int & kbx, + const int & iqs, + const uint32_t stride_col_y, + float * result) { + const block_ptq1_0 * bq = (const block_ptq1_0 *) vbq + kbx; + int sumi[ncols_dst][4] = {}; + + // Widen four bytes to 16-bit lanes so multiply-by-three cannot carry between bytes. +# pragma unroll + for (int g = 0; g < 4; ++g) { + const uint32_t packed = get_int_b4(bq->qs, g); + uint32_t v_lo = __byte_perm(packed, 0, 0x4140); + uint32_t v_hi = __byte_perm(packed, 0, 0x4342); + +# pragma unroll + for (int t = 0; t < 5; ++t) { + const uint32_t w_lo = v_lo * 3; + const uint32_t w_hi = v_hi * 3; + v_lo = w_lo & 0x00FF00FF; + v_hi = w_hi & 0x00FF00FF; + + const int q = __vsub4(__byte_perm(w_lo, w_hi, 0x7531), 0x01010101); + const int e = t * 16 + 4 * g; +# pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + const int u = get_int_b4(bq8_1[j * stride_col_y + iqs + (e >> 5)].qs, (e & 31) >> 2); + sumi[j][e >> 5] = ggml_cuda_dp4a(q, u, sumi[j][e >> 5]); + } + } + } + +# pragma unroll + for (int g = 0; g < 2; ++g) { + const uint32_t packed = get_int_b4(bq->qs + 16, g); + uint32_t v_lo = __byte_perm(packed, 0, 0x4140); + uint32_t v_hi = __byte_perm(packed, 0, 0x4342); + +# pragma unroll + for (int t = 0; t < 5; ++t) { + const uint32_t w_lo = v_lo * 3; + const uint32_t w_hi = v_hi * 3; + v_lo = w_lo & 0x00FF00FF; + v_hi = w_hi & 0x00FF00FF; + + const int q = __vsub4(__byte_perm(w_lo, w_hi, 0x7531), 0x01010101); + const int e = 80 + t * 8 + 4 * g; +# pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + const int u = get_int_b4(bq8_1[j * stride_col_y + iqs + (e >> 5)].qs, (e & 31) >> 2); + sumi[j][e >> 5] = ggml_cuda_dp4a(q, u, sumi[j][e >> 5]); + } + } + } + + uint32_t v = (uint32_t) bq->qh[0] | ((uint32_t) bq->qh[1] << 16); +# pragma unroll + for (int t = 0; t < 4; t += 2) { + const uint32_t w0 = v * 3; + v = w0 & 0x00FF00FF; + const uint32_t w1 = v * 3; + v = w1 & 0x00FF00FF; + + const int q = __vsub4(__byte_perm(w0, w1, 0x7531), 0x01010101); +# pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + const int u = get_int_b4(bq8_1[j * stride_col_y + iqs + 3].qs, 6 + t / 2); + sumi[j][3] = ggml_cuda_dp4a(q, u, sumi[j][3]); + } + } + + const float d = (float) bq->d; +# pragma unroll + for (int j = 0; j < ncols_dst; ++j) { + float acc = 0.0f; +# pragma unroll + for (int k = 0; k < 4; ++k) { + acc += __low2float(bq8_1[j * stride_col_y + iqs + k].ds) * (float) sumi[j][k]; + } + result[j] = d * acc; + } +} +#endif + +// PTQ1_0 x Q8_1. One call consumes the full 128-weight block. +static __device__ __forceinline__ float vec_dot_ptq1_0_q8_1(const void * __restrict__ vbq, + const block_q8_1 * __restrict__ bq8_1, + const int & kbx, + const int & iqs) { +#if defined(GGML_USE_HIP) + const block_ptq1_0 * bq = (const block_ptq1_0 *) vbq + kbx; + int sumi[4] = { 0, 0, 0, 0 }; + +# pragma unroll + for (int m = 0; m < 16; ++m) { + uint32_t v = bq->qs[m]; +# pragma unroll + for (int t = 0; t < 5; ++t) { + const uint32_t w = v * 3; + const int q = (int) (w >> 8) - 1; + v = w & 0xFF; + const int e = t * 16 + m; + sumi[e >> 5] += q * (int) bq8_1[iqs + (e >> 5)].qs[e & 31]; + } + } + +# pragma unroll + for (int m = 0; m < 8; ++m) { + uint32_t v = bq->qs[16 + m]; +# pragma unroll + for (int t = 0; t < 5; ++t) { + const uint32_t w = v * 3; + const int q = (int) (w >> 8) - 1; + v = w & 0xFF; + const int e = 80 + t * 8 + m; + sumi[e >> 5] += q * (int) bq8_1[iqs + (e >> 5)].qs[e & 31]; + } + } + +# pragma unroll + for (int h = 0; h < 2; ++h) { + uint32_t v = bq->qh[h]; +# pragma unroll + for (int t = 0; t < 4; ++t) { + const uint32_t w = v * 3; + const int q = (int) (w >> 8) - 1; + v = w & 0xFF; + const int e = 120 + t * 2 + h; + sumi[e >> 5] += q * (int) bq8_1[iqs + (e >> 5)].qs[e & 31]; + } + } + + float acc = 0.0f; +# pragma unroll + for (int k = 0; k < 4; ++k) { + acc += __low2float(bq8_1[iqs + k].ds) * (float) sumi[k]; + } + return (float) bq->d * acc; +#else + float result; + vec_dot_ptq1_0_q8_1_multi<1>(vbq, bq8_1, kbx, iqs, 0, &result); + return result; +#endif +} +static __device__ __forceinline__ float vec_dot_pq2_0_q8_1(const void * __restrict__ vbq, + const block_q8_1 * __restrict__ bq8_1, + const int & kbx, + const int & iqs) { const block_pq2_0 * bq2_0 = (const block_pq2_0 *) vbq + kbx; // Q2_0 group 128: 128 elements, ONE scale, processed as four 32-element chunks diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp index 4870bdcb968d..e85c76e0cca7 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.cpp +++ b/ggml/src/ggml-metal/ggml-metal-device.cpp @@ -871,6 +871,11 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta nsg = N_SG_PQ2_0; nr0 = N_R0_PQ2_0; } break; + case GGML_TYPE_PTQ1_0: + { + nsg = N_SG_PTQ1_0; + nr0 = N_R0_PTQ1_0; + } break; case GGML_TYPE_Q4_0: { nsg = N_SG_Q4_0; @@ -1110,6 +1115,11 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id(ggml_m nsg = N_SG_PQ2_0; nr0 = N_R0_PQ2_0; } break; + case GGML_TYPE_PTQ1_0: + { + nsg = N_SG_PTQ1_0; + nr0 = N_R0_PTQ1_0; + } break; case GGML_TYPE_Q4_0: { nsg = N_SG_Q4_0; diff --git a/ggml/src/ggml-metal/ggml-metal-device.m b/ggml/src/ggml-metal/ggml-metal-device.m index 08c0ef6ce67d..43b1d8c6f306 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.m +++ b/ggml/src/ggml-metal/ggml-metal-device.m @@ -1732,6 +1732,7 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -1762,6 +1763,7 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te case GGML_TYPE_Q1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_PQ2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: diff --git a/ggml/src/ggml-metal/ggml-metal-impl.h b/ggml/src/ggml-metal/ggml-metal-impl.h index 836cbbcb5c9e..1151e5099fc8 100644 --- a/ggml/src/ggml-metal/ggml-metal-impl.h +++ b/ggml/src/ggml-metal/ggml-metal-impl.h @@ -29,6 +29,8 @@ #define N_R0_PQ2_0 8 #define N_SG_PQ2_0 2 +#define N_R0_PTQ1_0 4 +#define N_SG_PTQ1_0 1 #define N_R0_Q4_0 4 #define N_SG_Q4_0 2 diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index 6be0f52015b2..38744c2d3cdd 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -2565,6 +2565,7 @@ int ggml_metal_op_mul_mat(ggml_metal_op_t ctx, int idx) { op->src[0]->type == GGML_TYPE_Q1_0 || op->src[0]->type == GGML_TYPE_Q2_0 || op->src[0]->type == GGML_TYPE_PQ2_0 || + op->src[0]->type == GGML_TYPE_PTQ1_0 || op->src[0]->type == GGML_TYPE_Q4_0 || op->src[0]->type == GGML_TYPE_Q4_1 || op->src[0]->type == GGML_TYPE_Q5_0 || diff --git a/ggml/src/ggml-metal/kernels/dequantize.h b/ggml/src/ggml-metal/kernels/dequantize.h index 1254d190c9b1..1e01dcd81d9e 100644 --- a/ggml/src/ggml-metal/kernels/dequantize.h +++ b/ggml/src/ggml-metal/kernels/dequantize.h @@ -161,6 +161,89 @@ void dequantize_pq2_0_t4(device const block_pq2_0 * xb, short il, thread type4 & reg = (type4) reg_f; } +// PTQ1_0: one lookup replaces the base-3 arithmetic. Entry b holds the five trits of +// byte b packed two bits each, low trit first, as 0/1/2 (subtract 1 for the value). +// Generated from, and verified against, the reference decode ((b*3^n & 0xFF)*3)>>8 for +// all 256 bytes and all five trit positions. This puts extraction on par with a 2-bit +// format: shift, mask, subtract, with a single cached table read amortised over 5 trits. +constant ushort ptq1_0_lut[256] = { + 0, 0, 256, 512, 64, 320, 576, 128, 384, 640, 16, 272, 528, 80, 336, 592, + 144, 400, 656, 32, 32, 288, 544, 96, 352, 608, 160, 416, 672, 4, 260, 516, + 68, 324, 580, 132, 388, 644, 20, 276, 276, 532, 84, 340, 596, 148, 404, 660, + 36, 292, 548, 100, 356, 612, 164, 420, 676, 8, 264, 520, 520, 72, 328, 584, + 136, 392, 648, 24, 280, 536, 88, 344, 600, 152, 408, 664, 40, 296, 552, 552, + 104, 360, 616, 168, 424, 680, 1, 257, 513, 65, 321, 577, 129, 385, 641, 17, + 273, 529, 81, 81, 337, 593, 145, 401, 657, 33, 289, 545, 97, 353, 609, 161, + 417, 673, 5, 261, 517, 69, 325, 325, 581, 133, 389, 645, 21, 277, 533, 85, + 341, 597, 149, 405, 661, 37, 293, 549, 101, 357, 357, 613, 165, 421, 677, 9, + 265, 521, 73, 329, 585, 137, 393, 649, 25, 281, 537, 89, 345, 601, 601, 153, + 409, 665, 41, 297, 553, 105, 361, 617, 169, 425, 681, 2, 258, 514, 66, 322, + 578, 130, 130, 386, 642, 18, 274, 530, 82, 338, 594, 146, 402, 658, 34, 290, + 546, 98, 354, 610, 162, 162, 418, 674, 6, 262, 518, 70, 326, 582, 134, 390, + 646, 22, 278, 534, 86, 342, 598, 150, 406, 406, 662, 38, 294, 550, 102, 358, + 614, 166, 422, 678, 10, 266, 522, 74, 330, 586, 138, 394, 650, 650, 26, 282, + 538, 90, 346, 602, 154, 410, 666, 42, 298, 554, 106, 362, 618, 170, 426, 682, +}; + +// PTQ1_0 stores trits base-3 packed, so element order is not positional: it follows +// the CPU codec's 16-then-8 byte staging over qs, then qh at 4 trits per byte. +// Map an element index to its byte and trit rather than assuming contiguity. +inline float ptq1_0_elem(device const block_ptq1_0 * xb, int e) { + uchar b; + short n; + if (e < 80) { // qs[0..15], chunk of 16, 5 trits per byte + b = xb->qs[e & 15]; n = e >> 4; + } else if (e < 120) { // qs[16..23], chunk of 8 + const int t = e - 80; + b = xb->qs[16 + (t & 7)]; n = t >> 3; + } else { // qh[0..1], 4 trits per byte + const int t = e - 120; + b = xb->qh[t & 1]; n = t >> 1; + } + return (float) ((int) ((ptq1_0_lut[b] >> (2*n)) & 3) - 1); +} + +template +void dequantize_ptq1_0(device const block_ptq1_0 * xb, short il, thread type4x4 & reg) { + const float d = xb->d; + + float4x4 reg_f; + + // il 0..4 addresses the 16-byte chunk at a single trit index, so the pow3 factor + // and the branch hoist out and this becomes 16 plain byte loads. il 5..7 straddle + // trit indices, or reach into qh, and take the general accessor. + if (il < 5) { + const short sh = 2*il; + device const uchar * qs = xb->qs; + for (int i = 0; i < 4; i++) { + for (int j = 0; j < 4; j++) { + reg_f[i][j] = (float) ((int) ((ptq1_0_lut[qs[i*4 + j]] >> sh) & 3) - 1) * d; + } + } + } else { + const int base = il * 16; + for (int i = 0; i < 4; i++) { + for (int j = 0; j < 4; j++) { + reg_f[i][j] = ptq1_0_elem(xb, base + i*4 + j) * d; + } + } + } + + reg = (type4x4) reg_f; +} + +template +void dequantize_ptq1_0_t4(device const block_ptq1_0 * xb, short il, thread type4 & reg) { + const float d = xb->d; + const int base = il * 4; + + float4 reg_f; + for (int j = 0; j < 4; j++) { + reg_f[j] = ptq1_0_elem(xb, base + j) * d; + } + reg = (type4) reg_f; +} + template void dequantize_q4_0(device const block_q4_0 * xb, short il, thread type4x4 & reg) { device const uint16_t * qs = ((device const uint16_t *)xb + 1); diff --git a/ggml/src/ggml-metal/kernels/mul_mm.metal b/ggml/src/ggml-metal/kernels/mul_mm.metal index 443f41d52332..b0029fdc6ad2 100644 --- a/ggml/src/ggml-metal/kernels/mul_mm.metal +++ b/ggml/src/ggml-metal/kernels/mul_mm.metal @@ -744,6 +744,7 @@ template [[host_name("kernel_mul_mm_bf16_f32")]] kernel mul_mm_t kernel_mul_m template [[host_name("kernel_mul_mm_q1_0_f32")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q2_0_f32")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_pq2_0_f32")]] kernel mul_mm_t kernel_mul_mm; +template [[host_name("kernel_mul_mm_ptq1_0_f32")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q4_0_f32")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q4_1_f32")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q5_0_f32")]] kernel mul_mm_t kernel_mul_mm; @@ -771,6 +772,7 @@ template [[host_name("kernel_mul_mm_f16_f16")]] kernel mul_mm_t kernel_mul_m template [[host_name("kernel_mul_mm_q1_0_f16")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q2_0_f16")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_pq2_0_f16")]] kernel mul_mm_t kernel_mul_mm; +template [[host_name("kernel_mul_mm_ptq1_0_f16")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q4_0_f16")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q4_1_f16")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q5_0_f16")]] kernel mul_mm_t kernel_mul_mm; @@ -807,6 +809,7 @@ template [[host_name("kernel_mul_mm_id_bf16_f32")]] kernel mul_mm_id kernel_m template [[host_name("kernel_mul_mm_id_q1_0_f32")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q2_0_f32")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_pq2_0_f32")]] kernel mul_mm_id kernel_mul_mm_id; +template [[host_name("kernel_mul_mm_id_ptq1_0_f32")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q4_0_f32")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q4_1_f32")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q5_0_f32")]] kernel mul_mm_id kernel_mul_mm_id; @@ -834,6 +837,7 @@ template [[host_name("kernel_mul_mm_id_f16_f16")]] kernel mul_mm_id kernel_m template [[host_name("kernel_mul_mm_id_q1_0_f16")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q2_0_f16")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_pq2_0_f16")]] kernel mul_mm_id kernel_mul_mm_id; +template [[host_name("kernel_mul_mm_id_ptq1_0_f16")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q4_0_f16")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q4_1_f16")]] kernel mul_mm_id kernel_mul_mm_id; template [[host_name("kernel_mul_mm_id_q5_0_f16")]] kernel mul_mm_id kernel_mul_mm_id; diff --git a/ggml/src/ggml-metal/kernels/mul_mv.metal b/ggml/src/ggml-metal/kernels/mul_mv.metal index bd4482ecf205..222902cc7e0b 100644 --- a/ggml/src/ggml-metal/kernels/mul_mv.metal +++ b/ggml/src/ggml-metal/kernels/mul_mv.metal @@ -510,6 +510,160 @@ kernel void kernel_mul_mv_q2_0_f32( kernel_mul_mv_q2_0_f32_impl(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg); } +// Byte-owning dot for PTQ1_0. The previous mapping gave each thread 16 contiguous +// weights, which in base-3 layout means every thread reloads the same qs bytes for a +// different trit index -- five times the byte traffic for the same 24 bytes. Here a +// thread instead owns whole bytes and consumes all five trits of each one, so a block's +// 26 bytes are loaded exactly once across the eight threads that cover it. The cost is +// that the y reads become strided rather than contiguous, which is the cheaper side: +// weight traffic is what decode is bound by. +// +// Byte ownership for thread it in 0..7, matching the CPU codec's element order: +// qs[2*it], qs[2*it+1] -> elements n*16 + m for n in 0..4 +// qs[16 + it] -> elements 80 + n*8 + it for n in 0..4 +// qh[it] (it < 2 only) -> elements 120 + n*2 + it for n in 0..3 +// Dot against y already staged in registers by the caller and reused across all nr0 +// rows. Trits come out of the byte by the base-3 remainder recurrence +// t = (v*3) >> 8, v = (v*3) & 0xFF, v starting at the byte +// which is two integer ops per trit with no table and no pow3 lookup. A 256-entry +// lookup was measurably worse here: it adds a dependent load per byte, and decode on +// this kernel is instruction-bound, not bandwidth-bound. Verified against the +// reference decode for all 256 bytes and all five positions. +// Accumulates the raw 0/1/2 trit and subtracts the staged y sum once, so there is no +// per-element offset correction. +inline float ptq1_0_dot_reg(device const block_ptq1_0 * qb, thread const float * yl, float sumy, short it) { + float acc = 0.f; + short c = 0; + + FOR_UNROLL (short k = 0; k < 2; ++k) { + ushort v = qb->qs[2*it + k]; + FOR_UNROLL (short n = 0; n < 5; ++n) { + const ushort w = v * 3; + acc += (float) (w >> 8) * yl[c++]; + v = w & 0xFF; + } + } + + { + ushort v = qb->qs[16 + it]; + FOR_UNROLL (short n = 0; n < 5; ++n) { + const ushort w = v * 3; + acc += (float) (w >> 8) * yl[c++]; + v = w & 0xFF; + } + } + + // qh holds 8 elements; give every thread exactly one so all eight do 16 elements. + // Element 120+it sits at trit it>>1 of byte qh[it&1], so step the recurrence to it. + { + ushort v = qb->qh[it & 1]; + const short n = it >> 1; + for (short i = 0; i < n; ++i) { + v = (v * 3) & 0xFF; + } + acc += (float) ((v * 3) >> 8) * yl[c++]; + } + + return (acc - sumy) * (float) qb->d; +} + +template +void kernel_mul_mv_ptq1_0_f32_impl( + args_t args, + device const char * src0, + device const char * src1, + device char * dst, + threadgroup char * shmem, + uint3 tgpig, + ushort tiisg, + ushort sgitg) { + const short NSG = FC_mul_mv_nsg; + + const int nb = args.ne00/QK_PTQ1_0; + + const int r0 = tgpig.x; + const int r1 = tgpig.y; + const int im = tgpig.z; + + const int first_row = (r0 * NSG + sgitg) * nr0; + + const uint i12 = im%FC_mul_mv_ne12; + const uint i13 = im/FC_mul_mv_ne12; + + const uint64_t offset1 = r1*args.nb11 + (i12)*args.nb12 + (i13)*args.nb13; + + device const float * y = (device const float *) (src1 + offset1); + + device const block_ptq1_0 * ax[nr0]; + for (int row = 0; row < nr0; ++row) { + const uint64_t offset0 = (first_row + row)*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03; + ax[row] = (device const block_ptq1_0 *) ((device char *) src0 + offset0); + } + + float yl[16]; + float sumf[nr0] = {0.f}; + + // eight threads cover one 128-weight block; each owns whole bytes, not a + // contiguous element span, so the block's bytes are read once in total + const short ix = (tiisg/8); + const short it = (tiisg%8); + + device const float * yb = y + ix*QK_PTQ1_0; + + for (int ib = ix; ib < nb; ib += N_SIMDWIDTH/8) { + // stage this thread's y values once, then reuse them for every row + float sumy = 0.f; + short c = 0; + + FOR_UNROLL (short k = 0; k < 2; ++k) { + const short m = 2*it + k; + FOR_UNROLL (short n = 0; n < 5; ++n) { + const float v = yb[n*16 + m]; + yl[c++] = v; + sumy += v; + } + } + FOR_UNROLL (short n = 0; n < 5; ++n) { + const float v = yb[80 + n*8 + it]; + yl[c++] = v; + sumy += v; + } + { + const float v = yb[120 + it]; + yl[c++] = v; + sumy += v; + } + + FOR_UNROLL (short row = 0; row < nr0; row++) { + sumf[row] += ptq1_0_dot_reg(ax[row] + ib, yl, sumy, it); + } + + yb += QK_PTQ1_0 * (N_SIMDWIDTH/8); + } + + device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0; + + for (int row = 0; row < nr0; ++row) { + const float tot = simd_sum(sumf[row]); + + if (tiisg == 0 && first_row + row < args.ne01) { + dst_f32[first_row + row] = tot; + } + } +} + +[[host_name("kernel_mul_mv_ptq1_0_f32")]] +kernel void kernel_mul_mv_ptq1_0_f32( + constant ggml_metal_kargs_mul_mv & args, + device const char * src0, + device const char * src1, + device char * dst, + uint3 tgpig[[threadgroup_position_in_grid]], + ushort tiisg[[thread_index_in_simdgroup]], + ushort sgitg[[simdgroup_index_in_threadgroup]]) { + kernel_mul_mv_ptq1_0_f32_impl(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg); +} + template void kernel_mul_mv_pq2_0_f32_impl( args_t args, @@ -995,6 +1149,10 @@ template [[host_name("kernel_mul_mv_ext_q2_0_f32_r1_5")]] kernel mul_mv_ext_q4 template [[host_name("kernel_mul_mv_ext_pq2_0_f32_r1_2")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<2, block_pq2_0, 128, dequantize_pq2_0_t4>; template [[host_name("kernel_mul_mv_ext_pq2_0_f32_r1_3")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<3, block_pq2_0, 128, dequantize_pq2_0_t4>; template [[host_name("kernel_mul_mv_ext_pq2_0_f32_r1_4")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<4, block_pq2_0, 128, dequantize_pq2_0_t4>; +template [[host_name("kernel_mul_mv_ext_ptq1_0_f32_r1_2")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<2, block_ptq1_0, 128, dequantize_ptq1_0_t4>; +template [[host_name("kernel_mul_mv_ext_ptq1_0_f32_r1_3")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<3, block_ptq1_0, 128, dequantize_ptq1_0_t4>; +template [[host_name("kernel_mul_mv_ext_ptq1_0_f32_r1_4")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<4, block_ptq1_0, 128, dequantize_ptq1_0_t4>; +template [[host_name("kernel_mul_mv_ext_ptq1_0_f32_r1_5")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<5, block_ptq1_0, 128, dequantize_ptq1_0_t4>; template [[host_name("kernel_mul_mv_ext_pq2_0_f32_r1_5")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<5, block_pq2_0, 128, dequantize_pq2_0_t4>; template [[host_name("kernel_mul_mv_ext_q4_0_f32_r1_2")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<2, block_q4_0, 32, dequantize_q4_0_t4>; diff --git a/ggml/src/ggml-metal/kernels/quantize.metal b/ggml/src/ggml-metal/kernels/quantize.metal index c3ac9eccdbfd..a49d8ad7e234 100644 --- a/ggml/src/ggml-metal/kernels/quantize.metal +++ b/ggml/src/ggml-metal/kernels/quantize.metal @@ -284,6 +284,7 @@ typedef decltype(kernel_get_rows_q) get_rows_q_t template [[host_name("kernel_get_rows_q1_0")]] kernel get_rows_q_t kernel_get_rows_q; template [[host_name("kernel_get_rows_q2_0")]] kernel get_rows_q_t kernel_get_rows_q; template [[host_name("kernel_get_rows_pq2_0")]] kernel get_rows_q_t kernel_get_rows_q; +template [[host_name("kernel_get_rows_ptq1_0")]] kernel get_rows_q_t kernel_get_rows_q; template [[host_name("kernel_get_rows_q4_0")]] kernel get_rows_q_t kernel_get_rows_q; template [[host_name("kernel_get_rows_q4_1")]] kernel get_rows_q_t kernel_get_rows_q; template [[host_name("kernel_get_rows_q5_0")]] kernel get_rows_q_t kernel_get_rows_q; diff --git a/ggml/src/ggml-quants.c b/ggml/src/ggml-quants.c index a5f8f9c352f3..bfb3892a599c 100644 --- a/ggml/src/ggml-quants.c +++ b/ggml/src/ggml-quants.c @@ -2195,6 +2195,107 @@ size_t quantize_pq2_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, return nrow * row_size; } +// ====================== PTQ1_0 (Prism ternary, group 128) ====================== +// Base-3 trit packing identical to upstream TQ1_0, but at block 128 so one fp16 +// scale covers 128 weights. qs is 24 bytes, which TQ1_0's fixed 32-then-16 byte +// staging cannot cover, so the stages are generalised to 32/16/8; at TQ1_0's +// 48-byte qs this reduces to exactly its original 32-then-16 behaviour. +static const size_t ptq1_0_stages[3] = {32, 16, 8}; + +void quantize_row_ptq1_0_ref(const float * GGML_RESTRICT x, block_ptq1_0 * GGML_RESTRICT y, int64_t k) { + assert(k % QK_PTQ1_0 == 0); + const int64_t nb = k / QK_PTQ1_0; + + for (int64_t i = 0; i < nb; i++) { + float amax = 0.0f; + for (int j = 0; j < QK_PTQ1_0; j++) { + amax = MAX(amax, fabsf(x[j])); + } + + const float d = amax; + const float id = d ? 1.0f/d : 0.0f; + + y[i].d = GGML_FP32_TO_FP16(d); + + size_t j = 0; + for (size_t s = 0; s < 3; ++s) { + const size_t c = ptq1_0_stages[s]; + for (; j + c <= sizeof(y->qs); j += c) { + for (size_t m = 0; m < c; ++m) { + uint8_t q = 0; + for (size_t n = 0; n < 5; ++n) { + int xi = lroundf(x[m + n*c] * id) + 1; // -1, 0, 1 -> 0, 1, 2 + q *= 3; + q += xi; + } + // ceiling division (243 == pow(3, 5)) + q = ((uint16_t)q * 256 + (243 - 1)) / 243; + y[i].qs[j + m] = q; + } + x += 5*c; + } + } + // 4 elements per byte + for (size_t h = 0; h < sizeof(y->qh); ++h) { + uint8_t q = 0; + for (size_t m = 0; m < 4; ++m) { + int xi = lroundf(x[h + m*sizeof(y->qh)] * id) + 1; + q *= 3; + q += xi; + } + // shift the first value to the most significant trit + q *= 3; + q = ((uint16_t)q * 256 + (243 - 1)) / 243; + y[i].qh[h] = q; + } + x += 4*sizeof(y->qh); + } +} + +void dequantize_row_ptq1_0(const block_ptq1_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + assert(k % QK_PTQ1_0 == 0); + const int64_t nb = k / QK_PTQ1_0; + + const uint8_t pow3[6] = {1, 3, 9, 27, 81, 243}; + + for (int64_t i = 0; i < nb; ++i) { + const float d = GGML_FP16_TO_FP32(x[i].d); + + size_t j = 0; + for (size_t s = 0; s < 3; ++s) { + const size_t c = ptq1_0_stages[s]; + for (; j + c <= sizeof(x->qs); j += c) { + for (size_t n = 0; n < 5; ++n) { + for (size_t m = 0; m < c; ++m) { + uint8_t q = x[i].qs[j + m] * pow3[n]; + int16_t xi = ((uint16_t) q * 3) >> 8; + *y++ = (float) (xi - 1) * d; + } + } + } + } + for (size_t n = 0; n < 4; ++n) { + for (size_t h = 0; h < sizeof(x->qh); ++h) { + uint8_t q = x[i].qh[h] * pow3[n]; + int16_t xi = ((uint16_t) q * 3) >> 8; + *y++ = (float) (xi - 1) * d; + } + } + } +} + +size_t quantize_ptq1_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { + (void)quant_weights; // ternary codes come from the weights themselves; an imatrix has no role + const size_t row_size = ggml_row_size(GGML_TYPE_PTQ1_0, n_per_row); + char * qrow = (char *)dst; + for (int64_t row = 0; row < nrow; ++row) { + quantize_row_ptq1_0_ref(src, (block_ptq1_0 *)qrow, n_per_row); + src += n_per_row; + qrow += row_size; + } + return nrow * row_size; +} + size_t quantize_q4_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { if (!quant_weights) { quantize_row_q4_0_ref(src, dst, (int64_t)nrow*n_per_row); @@ -5611,6 +5712,10 @@ bool ggml_validate_row_data(enum ggml_type type, const void * data, size_t nbyte { VALIDATE_ROW_DATA_D_F16_IMPL(block_pq2_0, data, nb); } break; + case GGML_TYPE_PTQ1_0: + { + VALIDATE_ROW_DATA_D_F16_IMPL(block_ptq1_0, data, nb); + } break; case GGML_TYPE_Q4_0: { VALIDATE_ROW_DATA_D_F16_IMPL(block_q4_0, data, nb); diff --git a/ggml/src/ggml-quants.h b/ggml/src/ggml-quants.h index 8a8370f48ef2..c9e5b2fe3587 100644 --- a/ggml/src/ggml-quants.h +++ b/ggml/src/ggml-quants.h @@ -17,6 +17,7 @@ extern "C" { GGML_API void quantize_row_q1_0_ref(const float * GGML_RESTRICT x, block_q1_0 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_q2_0_ref(const float * GGML_RESTRICT x, block_q2_0 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_pq2_0_ref(const float * GGML_RESTRICT x, block_pq2_0 * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_ptq1_0_ref(const float * GGML_RESTRICT x, block_ptq1_0 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_q4_0_ref(const float * GGML_RESTRICT x, block_q4_0 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_q4_1_ref(const float * GGML_RESTRICT x, block_q4_1 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_q5_0_ref(const float * GGML_RESTRICT x, block_q5_0 * GGML_RESTRICT y, int64_t k); @@ -47,6 +48,7 @@ GGML_API void quantize_row_iq2_s_ref (const float * GGML_RESTRICT x, block_iq2_ GGML_API void dequantize_row_q1_0(const block_q1_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_q2_0(const block_q2_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_pq2_0(const block_pq2_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_ptq1_0(const block_ptq1_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_q4_0(const block_q4_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_q4_1(const block_q4_1 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_q5_0(const block_q5_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); @@ -99,6 +101,7 @@ GGML_API size_t quantize_q6_K(const float * GGML_RESTRICT src, void * GGML_RESTR GGML_API size_t quantize_q1_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_q2_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_pq2_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); +GGML_API size_t quantize_ptq1_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_q4_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_q4_1(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_q5_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index 832e344ccffd..e30fb0d63329 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -4634,6 +4634,9 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { } #endif CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_Q1_0], matmul_q1_0_f16, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3) + // PTQ1_0 intentionally has no coopmat2 pipelines: dequant_funcs_cm2.glsl carries no + // PTQ1_0 decoder, so vulkan-shaders-gen skips cm2 generation for it and these symbols + // do not exist. It uses the scalar and coopmat1 matmul paths instead. CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_Q2_0], matmul_q2_0_f16, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3) CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_Q4_0], matmul_q4_0_f16, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3) CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_Q4_1], matmul_q4_1_f16, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3) @@ -4747,6 +4750,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { #endif CREATE_MM2(GGML_TYPE_Q1_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q1_0], matmul_q1_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, ); + CREATE_MM2(GGML_TYPE_PTQ1_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_PTQ1_0], matmul_ptq1_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, ); CREATE_MM2(GGML_TYPE_Q2_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q2_0], matmul_q2_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, ); CREATE_MM2(GGML_TYPE_Q4_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q4_0], matmul_q4_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, ); CREATE_MM2(GGML_TYPE_Q4_1, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q4_1], matmul_q4_1_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, ); @@ -4793,6 +4797,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { #endif CREATE_MM2(GGML_TYPE_Q1_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q1_0], matmul_id_subgroup_q1_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id); + CREATE_MM2(GGML_TYPE_PTQ1_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_PTQ1_0], matmul_id_subgroup_ptq1_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id); CREATE_MM2(GGML_TYPE_Q2_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_0], matmul_id_subgroup_q2_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id); CREATE_MM2(GGML_TYPE_Q4_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_0], matmul_id_subgroup_q4_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id); CREATE_MM2(GGML_TYPE_Q4_1, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_1], matmul_id_subgroup_q4_1_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id); @@ -4884,6 +4889,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { CREATE_MM_NODOT2(GGML_TYPE_BF16, pipeline_matmul_bf16, matmul_bf16, , wg_denoms, warptile, vk_mat_mat_push_constants, 3, , 0); CREATE_MM2(GGML_TYPE_Q1_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q1_0], matmul_q1_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0); + CREATE_MM2(GGML_TYPE_PTQ1_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_PTQ1_0], matmul_ptq1_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0); CREATE_MM2(GGML_TYPE_Q2_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q2_0], matmul_q2_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0); CREATE_MM2(GGML_TYPE_Q4_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q4_0], matmul_q4_0_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0); CREATE_MM2(GGML_TYPE_Q4_1, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q4_1], matmul_q4_1_f32, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0); @@ -4933,6 +4939,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { CREATE_MM2(GGML_TYPE_F16, pipeline_matmul_id_f16_f32, matmul_id_subgroup_f16_f32, wg_denoms, warptile_id, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size_16); CREATE_MM_NODOT2(GGML_TYPE_BF16, pipeline_matmul_id_bf16, matmul_id_subgroup_bf16, , wg_denoms, warptile_id, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size_16); CREATE_MM2(GGML_TYPE_Q1_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q1_0], matmul_id_subgroup_q1_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size); + CREATE_MM2(GGML_TYPE_PTQ1_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_PTQ1_0], matmul_id_subgroup_ptq1_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size); CREATE_MM2(GGML_TYPE_Q2_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_0], matmul_id_subgroup_q2_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size); CREATE_MM2(GGML_TYPE_Q4_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_0], matmul_id_subgroup_q4_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size); CREATE_MM2(GGML_TYPE_Q4_1, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_1], matmul_id_subgroup_q4_1_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size); @@ -4981,6 +4988,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { CREATE_MM2(GGML_TYPE_F16, pipeline_matmul_id_f16_f32, matmul_id_f16_f32, wg_denoms, warptile, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); CREATE_MM_NODOT2(GGML_TYPE_BF16, pipeline_matmul_id_bf16, matmul_id_bf16, , wg_denoms, warptile, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); CREATE_MM2(GGML_TYPE_Q1_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q1_0], matmul_id_q1_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); + CREATE_MM2(GGML_TYPE_PTQ1_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_PTQ1_0], matmul_id_ptq1_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); CREATE_MM2(GGML_TYPE_Q2_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_0], matmul_id_q2_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); CREATE_MM2(GGML_TYPE_Q4_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_0], matmul_id_q4_0_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); CREATE_MM2(GGML_TYPE_Q4_1, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_1], matmul_id_q4_1_f32, mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); @@ -5060,6 +5068,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { CREATE_MM(GGML_TYPE_BF16, pipeline_matmul_bf16, matmul_bf16, , wg_denoms, warptile, vk_mat_mat_push_constants, 3, , 0); CREATE_MM(GGML_TYPE_Q1_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q1_0].f32acc, matmul_q1_0_f32, , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0); + CREATE_MM(GGML_TYPE_PTQ1_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_PTQ1_0].f32acc, matmul_ptq1_0_f32, , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0); CREATE_MM(GGML_TYPE_Q2_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q2_0].f32acc, matmul_q2_0_f32, , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0); CREATE_MM(GGML_TYPE_Q4_0, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q4_0].f32acc, matmul_q4_0_f32, , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0); CREATE_MM(GGML_TYPE_Q4_1, pipeline_dequant_mul_mat_mat[GGML_TYPE_Q4_1].f32acc, matmul_q4_1_f32, , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0); @@ -5109,6 +5118,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { CREATE_MM(GGML_TYPE_BF16, pipeline_matmul_id_bf16, matmul_id_subgroup_bf16, , wg_denoms, warptile_id, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size_16); CREATE_MM(GGML_TYPE_Q1_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q1_0].f32acc, matmul_id_subgroup_q1_0_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size); + CREATE_MM(GGML_TYPE_PTQ1_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_PTQ1_0].f32acc, matmul_id_subgroup_ptq1_0_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size); CREATE_MM(GGML_TYPE_Q2_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_0].f32acc, matmul_id_subgroup_q2_0_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size); CREATE_MM(GGML_TYPE_Q4_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_0].f32acc, matmul_id_subgroup_q4_0_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size); CREATE_MM(GGML_TYPE_Q4_1, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_1].f32acc, matmul_id_subgroup_q4_1_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size); @@ -5139,6 +5149,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { CREATE_MM(GGML_TYPE_BF16, pipeline_matmul_id_bf16, matmul_id_bf16, , wg_denoms, warptile, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); CREATE_MM(GGML_TYPE_Q1_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q1_0].f32acc, matmul_id_q1_0_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); + CREATE_MM(GGML_TYPE_PTQ1_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_PTQ1_0].f32acc, matmul_id_ptq1_0_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); CREATE_MM(GGML_TYPE_Q2_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_0].f32acc, matmul_id_q2_0_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); CREATE_MM(GGML_TYPE_Q4_0, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_0].f32acc, matmul_id_q4_0_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); CREATE_MM(GGML_TYPE_Q4_1, pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_1].f32acc, matmul_id_q4_1_f32, , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0); @@ -5243,6 +5254,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_F16 ][i], "mul_mat_vec_f16_f32_f32", arr_dmmv_f16_f32_f32_len[reduc], arr_dmmv_f16_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2, 1, 1}, {wg_size_subgroup, 2, i+1}, 1, false, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_BF16][i], "mul_mat_vec_bf16_f32_f32", arr_dmmv_bf16_f32_f32_len[reduc], arr_dmmv_bf16_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2, 1, 1}, {wg_size_subgroup, 2, i+1}, 1, false, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_Q1_0][i], "mul_mat_vec_q1_0_f32_f32", arr_dmmv_q1_0_f32_f32_len[reduc], arr_dmmv_q1_0_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_PTQ1_0][i], "mul_mat_vec_ptq1_0_f32_f32", arr_dmmv_ptq1_0_f32_f32_len[reduc], arr_dmmv_ptq1_0_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_Q2_0][i], "mul_mat_vec_q2_0_f32_f32", arr_dmmv_q2_0_f32_f32_len[reduc], arr_dmmv_q2_0_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_Q4_0][i], "mul_mat_vec_q4_0_f32_f32", arr_dmmv_q4_0_f32_f32_len[reduc], arr_dmmv_q4_0_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_Q4_1][i], "mul_mat_vec_q4_1_f32_f32", arr_dmmv_q4_1_f32_f32_len[reduc], arr_dmmv_q4_1_f32_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size); @@ -5271,6 +5283,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_F16 ][i], "mul_mat_vec_f16_f16_f32", arr_dmmv_f16_f16_f32_len[reduc], arr_dmmv_f16_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2, 1, 1}, {wg_size_subgroup, 2, i+1}, 1, false, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_BF16][i], "mul_mat_vec_bf16_f16_f32", arr_dmmv_bf16_f16_f32_len[reduc], arr_dmmv_bf16_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2, 1, 1}, {wg_size_subgroup, 2, i+1}, 1, false, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_Q1_0][i], "mul_mat_vec_q1_0_f16_f32", arr_dmmv_q1_0_f16_f32_len[reduc], arr_dmmv_q1_0_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_PTQ1_0][i], "mul_mat_vec_ptq1_0_f16_f32", arr_dmmv_ptq1_0_f16_f32_len[reduc], arr_dmmv_ptq1_0_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_Q2_0][i], "mul_mat_vec_q2_0_f16_f32", arr_dmmv_q2_0_f16_f32_len[reduc], arr_dmmv_q2_0_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_Q4_0][i], "mul_mat_vec_q4_0_f16_f32", arr_dmmv_q4_0_f16_f32_len[reduc], arr_dmmv_q4_0_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_Q4_1][i], "mul_mat_vec_q4_1_f16_f32", arr_dmmv_q4_1_f16_f32_len[reduc], arr_dmmv_q4_1_f16_f32_data[reduc], "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq, i+1}, 1, true, use_subgroups, force_subgroup_size); @@ -5326,6 +5339,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_F16 ], "mul_mat_vec_id_f16_f32", arr_dmmv_id_f16_f32_f32_len[reduc], arr_dmmv_id_f16_f32_f32_data[reduc], "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {2, 1, 1}, {wg_size_subgroup, 2}, 1, false, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_BF16], "mul_mat_vec_id_bf16_f32", arr_dmmv_id_bf16_f32_f32_len[reduc], arr_dmmv_id_bf16_f32_f32_data[reduc], "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {2, 1, 1}, {wg_size_subgroup, 2}, 1, false, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_Q1_0], "mul_mat_vec_id_q1_0_f32", arr_dmmv_id_q1_0_f32_f32_len[reduc], arr_dmmv_id_q1_0_f32_f32_data[reduc], "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq}, 1, true, use_subgroups, force_subgroup_size); + ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_PTQ1_0], "mul_mat_vec_id_ptq1_0_f32", arr_dmmv_id_ptq1_0_f32_f32_len[reduc], arr_dmmv_id_ptq1_0_f32_f32_data[reduc], "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq}, 1, true, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_Q2_0], "mul_mat_vec_id_q2_0_f32", arr_dmmv_id_q2_0_f32_f32_len[reduc], arr_dmmv_id_q2_0_f32_f32_data[reduc], "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq}, 1, true, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_Q4_0], "mul_mat_vec_id_q4_0_f32", arr_dmmv_id_q4_0_f32_f32_len[reduc], arr_dmmv_id_q4_0_f32_f32_data[reduc], "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq}, 1, true, use_subgroups, force_subgroup_size); ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_Q4_1], "mul_mat_vec_id_q4_1_f32", arr_dmmv_id_q4_1_f32_f32_len[reduc], arr_dmmv_id_q4_1_f32_f32_data[reduc], "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {2*rm_stdq, 1, 1}, {wg_size_subgroup, 2*rm_stdq}, 1, true, use_subgroups, force_subgroup_size); @@ -5388,6 +5402,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { // dequant shaders ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_F32 ], "f32_to_f16", dequant_f32_len, dequant_f32_data, "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q1_0], "dequant_q1_0", dequant_q1_0_len, dequant_q1_0_data, "main", 2, 5 * sizeof(uint32_t), {256 * 8, 1, 1}, {}, 1); + ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_PTQ1_0], "dequant_ptq1_0", dequant_ptq1_0_len, dequant_ptq1_0_data, "main", 2, 5 * sizeof(uint32_t), {256 * 8, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q2_0], "dequant_q2_0", dequant_q2_0_len, dequant_q2_0_data, "main", 2, 5 * sizeof(uint32_t), {256 * 4, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q4_0], "dequant_q4_0", dequant_q4_0_len, dequant_q4_0_data, "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_Q4_1], "dequant_q4_1", dequant_q4_1_len, dequant_q4_1_data, "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1); @@ -5418,6 +5433,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_F16 ], "get_rows_f16", get_rows_f16_len, get_rows_f16_data, "main", 3, sizeof(vk_op_binary_push_constants), { 512, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_BF16], "get_rows_bf16", get_rows_bf16_len, get_rows_bf16_data, "main", 3, sizeof(vk_op_binary_push_constants), { 512, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_Q1_0], "get_rows_q1_0", get_rows_q1_0_len, get_rows_q1_0_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1); + ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_PTQ1_0], "get_rows_ptq1_0", get_rows_ptq1_0_len, get_rows_ptq1_0_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_Q2_0], "get_rows_q2_0", get_rows_q2_0_len, get_rows_q2_0_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_Q4_0], "get_rows_q4_0", get_rows_q4_0_len, get_rows_q4_0_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_Q4_1], "get_rows_q4_1", get_rows_q4_1_len, get_rows_q4_1_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1); @@ -5447,6 +5463,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_F16 ], "get_rows_f16_f32", get_rows_f16_f32_len, get_rows_f16_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), { 512, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_BF16], "get_rows_bf16_f32", get_rows_bf16_f32_len, get_rows_bf16_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), { 512, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_Q1_0], "get_rows_q1_0_f32", get_rows_q1_0_f32_len, get_rows_q1_0_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1); + ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_PTQ1_0], "get_rows_ptq1_0_f32", get_rows_ptq1_0_f32_len, get_rows_ptq1_0_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_Q2_0], "get_rows_q2_0_f32", get_rows_q2_0_f32_len, get_rows_q2_0_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_Q4_0], "get_rows_q4_0_f32", get_rows_q4_0_f32_len, get_rows_q4_0_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1); ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_Q4_1], "get_rows_q4_1_f32", get_rows_q4_1_f32_len, get_rows_q4_1_f32_data, "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1); @@ -7656,6 +7673,7 @@ static vk_pipeline ggml_vk_get_to_fp16(ggml_backend_vk_context * ctx, ggml_type switch (type) { case GGML_TYPE_F32: case GGML_TYPE_Q1_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: @@ -7731,6 +7749,7 @@ static vk_matmul_pipeline ggml_vk_get_mul_mat_mat_pipeline(ggml_backend_vk_conte switch (src0_type) { case GGML_TYPE_Q1_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: @@ -7761,7 +7780,17 @@ static vk_matmul_pipeline ggml_vk_get_mul_mat_mat_pipeline(ggml_backend_vk_conte if (ctx->device->coopmat2) { assert(src1_type == GGML_TYPE_F16); - return prec == GGML_PREC_DEFAULT ? ctx->device->pipeline_dequant_mul_mat_mat_f16[src0_type].f16acc : ctx->device->pipeline_dequant_mul_mat_mat_f16[src0_type].f32acc; + vk_matmul_pipeline2 & mmp = ctx->device->pipeline_dequant_mul_mat_mat_f16[src0_type]; + vk_matmul_pipeline pipelines = prec == GGML_PREC_DEFAULT ? mmp.f16acc : mmp.f32acc; + + // A type with no coopmat2 decoder (PTQ1_0) leaves these slots empty. The slots are + // always allocated, so returning one would hand back a pipeline whose l/m/s are null; + // nullptr instead routes the caller to its dequant + f16 matmul fallback. + if (pipelines->is_empty()) { + return nullptr; + } + + return pipelines; } if (ctx->device->coopmat_support) { return (ctx->device->fp16 && ctx->device->coopmat_acc_f16_support && prec == GGML_PREC_DEFAULT) ? ctx->device->pipeline_dequant_mul_mat_mat[src0_type].f16acc : ctx->device->pipeline_dequant_mul_mat_mat[src0_type].f32acc; @@ -7801,6 +7830,7 @@ static vk_pipeline ggml_vk_get_dequantize_mul_mat_vec(ggml_backend_vk_context * case GGML_TYPE_F16: case GGML_TYPE_BF16: case GGML_TYPE_Q1_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: @@ -7895,6 +7925,7 @@ static vk_matmul_pipeline ggml_vk_get_mul_mat_mat_id_pipeline(ggml_backend_vk_co switch (src0_type) { case GGML_TYPE_Q1_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: @@ -7929,6 +7960,11 @@ static vk_matmul_pipeline ggml_vk_get_mul_mat_mat_id_pipeline(ggml_backend_vk_co bool support_fp16acc = !mmp.f16acc->is_empty(); bool support_fp32acc = !mmp.f32acc->is_empty(); + if (!support_fp16acc && !support_fp32acc) { + // No coopmat2 decoder for this type; let the caller dequantize to f16 instead of asserting. + return nullptr; + } + if (support_fp16acc && (prefer_fp16acc || !support_fp32acc)) { return mmp.f16acc; } else { @@ -7968,6 +8004,7 @@ static vk_pipeline ggml_vk_get_dequantize_mul_mat_vec_id(ggml_backend_vk_context case GGML_TYPE_F16: case GGML_TYPE_BF16: case GGML_TYPE_Q1_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: @@ -18148,6 +18185,7 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm case GGML_TYPE_F16: case GGML_TYPE_BF16: case GGML_TYPE_Q1_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: @@ -18254,6 +18292,7 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm case GGML_TYPE_F16: case GGML_TYPE_BF16: case GGML_TYPE_Q1_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_Q2_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/dequant_funcs.glsl b/ggml/src/ggml-vulkan/vulkan-shaders/dequant_funcs.glsl index 627932bd3547..88be1360403a 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/dequant_funcs.glsl +++ b/ggml/src/ggml-vulkan/vulkan-shaders/dequant_funcs.glsl @@ -126,6 +126,20 @@ vec4 dequantize4(uint ib, uint iqs, uint a_offset) { } #endif +#if defined(DATA_A_PTQ1_0) +#include "ptq1_0.glsl" + +vec2 dequantize(uint ib, uint iqs, uint a_offset) { + return vec2(ptq1_0_trit(ib, a_offset, iqs), ptq1_0_trit(ib, a_offset, iqs + 1u)); +} +vec4 dequantize4(uint ib, uint iqs, uint a_offset) { + return vec4(ptq1_0_trit(ib, a_offset, iqs), + ptq1_0_trit(ib, a_offset, iqs + 1u), + ptq1_0_trit(ib, a_offset, iqs + 2u), + ptq1_0_trit(ib, a_offset, iqs + 3u)); +} +#endif + #if defined(DATA_A_Q1_0) vec2 dequantize(uint ib, uint iqs, uint a_offset) { const uint bits = uint(data_a[a_offset + ib].qs[iqs / 8u]) >> (iqs % 8u); @@ -564,6 +578,13 @@ vec2 get_dm(uint ib, uint a_offset) { } #endif +#if defined(DATA_A_PTQ1_0) +vec2 get_dm(uint ib, uint a_offset) { + const float d = float(data_a[a_offset + ib].d); + return vec2(d, 0); +} +#endif + #if defined(DATA_A_Q1_0) vec2 get_dm(uint ib, uint a_offset) { const float d = float(data_a[a_offset + ib].d); diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/dequant_ptq1_0.comp b/ggml/src/ggml-vulkan/vulkan-shaders/dequant_ptq1_0.comp new file mode 100644 index 000000000000..53baf41f65b9 --- /dev/null +++ b/ggml/src/ggml-vulkan/vulkan-shaders/dequant_ptq1_0.comp @@ -0,0 +1,54 @@ +#version 450 + +#include "dequant_head.glsl" + +layout(local_size_x = 256, local_size_y = 1, local_size_z = 1) in; + +layout (binding = 0) readonly buffer A {block_ptq1_0 data_a[];}; +layout (binding = 1) writeonly buffer D {D_TYPE data_b[];}; + +// One invocation per 8 elements: 16 invocations cover a 128-weight block. Element +// order follows the CPU codec (16-byte qs chunk, 8-byte qs chunk, then qh), so the +// index maths lives in one place rather than being reproduced per shader. +float ptq1_0_trit_l(uint ib, uint e) { + uint b; + uint n; + if (e < 80u) { + b = uint(data_a[ib].qs[e & 15u]); + n = e >> 4u; + } else if (e < 120u) { + const uint t = e - 80u; + b = uint(data_a[ib].qs[16u + (t & 7u)]); + n = t >> 3u; + } else { + const uint t = e - 120u; + b = uint(data_a[ib].qh[t & 1u]); + n = t >> 1u; + } + + uint v = b; + for (uint i = 0u; i < n; ++i) { + v = (v * 3u) & 0xFFu; + } + return float(int((v * 3u) >> 8u) - 1); +} + +void main() { + const uint i = gl_WorkGroupID.x * 4 + gl_LocalInvocationID.x / 64; + + const uint tid = gl_LocalInvocationID.x % 64; + const uint il = tid / 4; + const uint ir = tid % 4; + const uint ib = 4*i + ir; + if (ib >= p.nel / 128) { + return; + } + + const uint b_idx = 512*i + 128*ir + 8*il; + + const float d = float(data_a[ib].d); + + [[unroll]] for (uint l = 0; l < 8; ++l) { + data_b[b_idx + l] = D_TYPE(ptq1_0_trit_l(ib, 8*il + l) * d); + } +} diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/mul_mm.comp b/ggml/src/ggml-vulkan/vulkan-shaders/mul_mm.comp index 3df88044a5ee..9804c51a3f5a 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/mul_mm.comp +++ b/ggml/src/ggml-vulkan/vulkan-shaders/mul_mm.comp @@ -136,6 +136,9 @@ shared FLOAT_TYPEV2 buf_b[BN * SHMEM_STRIDE]; shared ACC_TYPE coopmat_stage[TM * TN * NUM_WARPS]; #endif +#if defined(DATA_A_PTQ1_0) +#include "ptq1_0.glsl" +#endif #include "mul_mm_id_funcs.glsl" #include "mul_mm_funcs.glsl" diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/mul_mm_funcs.glsl b/ggml/src/ggml-vulkan/vulkan-shaders/mul_mm_funcs.glsl index 7d852dced8ab..3125a3fcfa93 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/mul_mm_funcs.glsl +++ b/ggml/src/ggml-vulkan/vulkan-shaders/mul_mm_funcs.glsl @@ -152,6 +152,21 @@ void load_a_to_shmem(const uint pos_a, const uint row, const uint col, const uin const uint k_pair = row * LOAD_VEC_A / 2; store_a(col, k_pair, FLOAT_TYPEV2(v.xy)); store_a(col, k_pair + 1, FLOAT_TYPEV2(v.zw)); +#elif defined(DATA_A_PTQ1_0) + const uint idx = pos_a + col * p.stride_a / LOAD_VEC_A + row; + + const uint ib = idx / 16; + const uint grp = idx & 0xfu; // which 8-element group inside the block + const uint e0 = grp * 8u; + + const float d = float(data_a[ib].d); + + const uint k_pair = row * LOAD_VEC_A / 2; + [[unroll]] for (uint l = 0; l < 4; ++l) { + store_a(col, k_pair + l, FLOAT_TYPEV2( + ptq1_0_trit(ib, 0u, e0 + 2u*l) * d, + ptq1_0_trit(ib, 0u, e0 + 2u*l + 1u) * d)); + } #elif defined(DATA_A_Q1_0) const uint idx = pos_a + col * p.stride_a / LOAD_VEC_A + row; diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/ptq1_0.glsl b/ggml/src/ggml-vulkan/vulkan-shaders/ptq1_0.glsl new file mode 100644 index 000000000000..7eedd832b14f --- /dev/null +++ b/ggml/src/ggml-vulkan/vulkan-shaders/ptq1_0.glsl @@ -0,0 +1,38 @@ +#ifndef PTQ1_0_GLSL +#define PTQ1_0_GLSL + +// Shared PTQ1_0 trit accessor. Lives in its own header because two consumers need it +// from different include chains: dequant_funcs.glsl (mul_mat_vec, get_rows, +// copy_from_quant) and mul_mm.comp (via mul_mm_funcs.glsl), and mul_mm.comp does not +// include dequant_funcs.glsl. Duplicating it would leave two copies that must stay in +// step with the CPU codec in ggml-quants.c, where a divergence shows up as wrong +// matmul results rather than a build error. +// +// Element order is not positional: a 16-byte chunk of qs where byte j carries +// elements t*16+j, then an 8-byte chunk carrying 80 + t*8 + (j-16), then qh at four +// trits per byte carrying 120 + t*2 + h. Trits come out by the base-3 remainder +// recurrence t = (v*3)>>8, v = (v*3)&0xFF. +float ptq1_0_trit(uint ib, uint a_offset, uint e) { + uint b; + uint n; + if (e < 80u) { + b = uint(data_a[a_offset + ib].qs[e & 15u]); + n = e >> 4u; + } else if (e < 120u) { + const uint t = e - 80u; + b = uint(data_a[a_offset + ib].qs[16u + (t & 7u)]); + n = t >> 3u; + } else { + const uint t = e - 120u; + b = uint(data_a[a_offset + ib].qh[t & 1u]); + n = t >> 1u; + } + + uint v = b; + for (uint i = 0u; i < n; ++i) { + v = (v * 3u) & 0xFFu; + } + return float(int((v * 3u) >> 8u) - 1); +} + +#endif // PTQ1_0_GLSL diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/types.glsl b/ggml/src/ggml-vulkan/vulkan-shaders/types.glsl index adb1bb8b32b5..469aad51ce89 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/types.glsl +++ b/ggml/src/ggml-vulkan/vulkan-shaders/types.glsl @@ -211,6 +211,27 @@ struct block_q1_0 #define A_TYPE block_q1_0 #endif +// PTQ1_0: ternary at group 128, base-3 packed five trits per byte. +// Field order mirrors block_ptq1_0 in ggml-common.h EXACTLY -- qs, then qh, then d. +// Unlike q1_0 the scale is LAST, and getting that wrong silently misindexes every +// block rather than failing loudly. +#define QUANT_K_PTQ1_0 128 +#define QUANT_R_PTQ1_0 1 + +struct block_ptq1_0 +{ + uint8_t qs[24]; + uint8_t qh[2]; + float16_t d; +}; + +#if defined(DATA_A_PTQ1_0) +#define QUANT_K QUANT_K_PTQ1_0 +#define QUANT_R QUANT_R_PTQ1_0 +#define QUANT_AUXF 1 +#define A_TYPE block_ptq1_0 +#endif + #define QUANT_K_Q2_0 64 #define QUANT_R_Q2_0 1 diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp index 770496b112c6..83af8af068fe 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp +++ b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp @@ -50,6 +50,7 @@ const std::vector type_names = { "f32", "f16", "q1_0", + "ptq1_0", "q2_0", "q4_0", "q4_1", @@ -583,7 +584,7 @@ void matmul_shaders(bool fp16, MatMulIdType matmul_id_type, bool coopmat, bool c for (const auto& tname : type_names) { std::string load_vec_quant = "2"; - if ((tname == "q1_0") || (tname == "q4_0") || (tname == "q4_1") || (tname == "q5_1") || (tname == "iq1_s") || (tname == "iq1_m") || (tname == "iq2_xxs") || (tname == "iq2_xs") || (tname == "iq2_s")) + if ((tname == "q1_0") || (tname == "ptq1_0") || (tname == "q4_0") || (tname == "q4_1") || (tname == "q5_1") || (tname == "iq1_s") || (tname == "iq1_m") || (tname == "iq2_xxs") || (tname == "iq2_xs") || (tname == "iq2_s")) load_vec_quant = "8"; else if ((tname == "q2_0") || (tname == "q5_0") || (tname == "q8_0") || (tname == "q2_k") || (tname == "q4_k") || (tname == "q5_k") || (tname == "iq3_xxs") || (tname == "iq3_s") || (tname == "iq4_xs") || (tname == "iq4_nl") || (tname == "mxfp4") || (tname == "nvfp4")) load_vec_quant = "4"; @@ -591,6 +592,13 @@ void matmul_shaders(bool fp16, MatMulIdType matmul_id_type, bool coopmat, bool c if (tname == "bf16") { continue; } + // PTQ1_0 has no coopmat2 decoder: dequant_funcs_cm2.glsl carries no PTQ1_0 entry, + // so emitting mul_mm_cm2 for it fails shader compilation and takes the whole + // Vulkan build down, not just this type. Skip it; it falls back to the scalar and + // coopmat1 matmul paths, which are the ones implemented and tested. + if (coopmat2 && tname == "ptq1_0") { + continue; + } std::string data_a_key = "DATA_A_" + to_uppercase(tname); // For aligned matmul loads diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index 6849721d8c60..2e7416070d04 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -697,6 +697,14 @@ static const struct ggml_type_traits type_traits[GGML_TYPE_COUNT] = { .to_float = (ggml_to_float_t) dequantize_row_pq2_0, .from_float_ref = (ggml_from_float_t) quantize_row_pq2_0_ref, }, + [GGML_TYPE_PTQ1_0] = { + .type_name = "ptq1_0", + .blck_size = QK_PTQ1_0, + .type_size = sizeof(block_ptq1_0), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_ptq1_0, + .from_float_ref = (ggml_from_float_t) quantize_row_ptq1_0_ref, + }, [GGML_TYPE_Q4_0] = { .type_name = "q4_0", .blck_size = QK4_0, @@ -1443,6 +1451,7 @@ enum ggml_type ggml_ftype_to_ggml_type(enum ggml_ftype ftype) { case GGML_FTYPE_MOSTLY_Q1_0: wtype = GGML_TYPE_Q1_0; break; case GGML_FTYPE_MOSTLY_Q2_0: wtype = GGML_TYPE_Q2_0; break; case GGML_FTYPE_MOSTLY_PQ2_0: wtype = GGML_TYPE_PQ2_0; break; + case GGML_FTYPE_MOSTLY_PTQ1_0: wtype = GGML_TYPE_PTQ1_0; break; case GGML_FTYPE_MOSTLY_Q5_0: wtype = GGML_TYPE_Q5_0; break; case GGML_FTYPE_MOSTLY_Q5_1: wtype = GGML_TYPE_Q5_1; break; case GGML_FTYPE_MOSTLY_Q8_0: wtype = GGML_TYPE_Q8_0; break; @@ -8050,6 +8059,7 @@ size_t ggml_quantize_chunk( case GGML_TYPE_Q1_0: result = quantize_q1_0 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q2_0: result = quantize_q2_0 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_PQ2_0: result = quantize_pq2_0(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_PTQ1_0: result = quantize_ptq1_0(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q4_0: result = quantize_q4_0 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q4_1: result = quantize_q4_1 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q5_0: result = quantize_q5_0 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; diff --git a/gguf-py/gguf/constants.py b/gguf-py/gguf/constants.py index 7777e7efa94d..d2284150ab6b 100644 --- a/gguf-py/gguf/constants.py +++ b/gguf-py/gguf/constants.py @@ -5384,7 +5384,8 @@ class GGMLQuantizationType(IntEnum): NVFP4 = 40 Q1_0 = 41 Q2_0 = 42 - PQ2_0 = 142 + PQ2_0 = 142 + PTQ1_0 = 143 class ExpertGatingFuncType(IntEnum): @@ -5441,7 +5442,8 @@ class LlamaFileType(IntEnum): MOSTLY_NVFP4 = 39 # except 1d tensors MOSTLY_Q1_0 = 40 # except 1d tensors MOSTLY_Q2_0 = 41 # except 1d tensors - MOSTLY_PQ2_0 = 128 # except 1d tensors + MOSTLY_PQ2_0 = 128 # except 1d tensors + MOSTLY_PTQ1_0 = 129 # except 1d tensors GUESSED = 1024 # not specified in the model file @@ -5578,7 +5580,8 @@ class VisionProjectorType: GGMLQuantizationType.NVFP4: (64, 4 + 32), GGMLQuantizationType.Q1_0: (128, 2 + 16), GGMLQuantizationType.Q2_0: (64, 2 + 16), - GGMLQuantizationType.PQ2_0: (128, 2 + 32), + GGMLQuantizationType.PQ2_0: (128, 2 + 32), + GGMLQuantizationType.PTQ1_0: (128, 2 + 24 + 2), } diff --git a/include/llama.h b/include/llama.h index e424d5c40100..d1aba602816a 100644 --- a/include/llama.h +++ b/include/llama.h @@ -158,6 +158,7 @@ extern "C" { LLAMA_FTYPE_MOSTLY_Q2_0 = 41, // except 1d tensors LLAMA_FTYPE_MOSTLY_PQ2_0 = 141, // except 1d tensors (Prism group-128 Q2_0; matches published PQ2_0 ggufs) LLAMA_FTYPE_MOSTLY_PQ2_0_LEGACY = 142, // pre-rename value for the same format, still found in published ggufs + LLAMA_FTYPE_MOSTLY_PTQ1_0 = 143, // except 1d tensors (Prism group-128 ternary, 1.75 bpw) LLAMA_FTYPE_GUESSED = 1024, // not specified in the model file }; diff --git a/src/llama-model-loader.cpp b/src/llama-model-loader.cpp index c43499aa234a..835d172d33f7 100644 --- a/src/llama-model-loader.cpp +++ b/src/llama-model-loader.cpp @@ -40,6 +40,7 @@ const char * llama_ftype_name(llama_ftype ftype) { case LLAMA_FTYPE_MOSTLY_Q1_0: name = LLAMA_FTYPE_PREFIX "Q1_0"; break; case LLAMA_FTYPE_MOSTLY_Q2_0: name = LLAMA_FTYPE_PREFIX "Q2_0"; break; case LLAMA_FTYPE_MOSTLY_PQ2_0: name = LLAMA_FTYPE_PREFIX "PQ2_0 - 2.13 bpw (group 128)"; break; + case LLAMA_FTYPE_MOSTLY_PTQ1_0: name = LLAMA_FTYPE_PREFIX "PTQ1_0 - 1.75 bpw ternary (group 128)"; break; // ggufs packed before the Q2_0_G128 -> PQ2_0 rename carry the old ftype value. // They load and compute correctly; name it so it does not report as unknown. case LLAMA_FTYPE_MOSTLY_PQ2_0_LEGACY: name = LLAMA_FTYPE_PREFIX "PQ2_0 - 2.13 bpw (group 128, legacy ftype)"; break; @@ -775,6 +776,7 @@ llama_model_loader::llama_model_loader( case GGML_TYPE_Q1_0: ftype = LLAMA_FTYPE_MOSTLY_Q1_0; break; case GGML_TYPE_Q2_0: ftype = LLAMA_FTYPE_MOSTLY_Q2_0; break; case GGML_TYPE_PQ2_0: ftype = LLAMA_FTYPE_MOSTLY_PQ2_0; break; + case GGML_TYPE_PTQ1_0: ftype = LLAMA_FTYPE_MOSTLY_PTQ1_0; break; default: { LLAMA_LOG_WARN("%s: unknown type %s\n", __func__, ggml_type_name(type_max)); diff --git a/src/llama-quant.cpp b/src/llama-quant.cpp index a06b44929c40..121e610abf5a 100644 --- a/src/llama-quant.cpp +++ b/src/llama-quant.cpp @@ -393,6 +393,7 @@ static ggml_type tensor_type_fallback(quantize_state_impl & qs, const ggml_tenso case GGML_TYPE_IQ3_S: // types on the right: block size 32 case GGML_TYPE_IQ4_XS: return_type = GGML_TYPE_IQ4_NL; break; case GGML_TYPE_Q2_0: + case GGML_TYPE_PTQ1_0: case GGML_TYPE_PQ2_0: case GGML_TYPE_Q2_K: case GGML_TYPE_Q3_K: @@ -499,7 +500,7 @@ static ggml_type llama_tensor_get_type_impl(quantize_state_impl & qs, ggml_type else if (ftype == LLAMA_FTYPE_MOSTLY_IQ3_XXS) { new_type = GGML_TYPE_IQ3_S; } - else if (ftype == LLAMA_FTYPE_MOSTLY_TQ1_0 || ftype == LLAMA_FTYPE_MOSTLY_TQ2_0 || ftype == LLAMA_FTYPE_MOSTLY_Q2_0 || ftype == LLAMA_FTYPE_MOSTLY_PQ2_0) { + else if (ftype == LLAMA_FTYPE_MOSTLY_TQ1_0 || ftype == LLAMA_FTYPE_MOSTLY_TQ2_0 || ftype == LLAMA_FTYPE_MOSTLY_Q2_0 || ftype == LLAMA_FTYPE_MOSTLY_PQ2_0 || ftype == LLAMA_FTYPE_MOSTLY_PTQ1_0) { new_type = GGML_TYPE_Q4_K; } } @@ -821,6 +822,7 @@ ggml_type llama_ftype_get_default_type(llama_ftype ftype) { case LLAMA_FTYPE_MOSTLY_Q1_0: return GGML_TYPE_Q1_0; case LLAMA_FTYPE_MOSTLY_Q2_0: return GGML_TYPE_Q2_0; case LLAMA_FTYPE_MOSTLY_PQ2_0: return GGML_TYPE_PQ2_0; + case LLAMA_FTYPE_MOSTLY_PTQ1_0: return GGML_TYPE_PTQ1_0; case LLAMA_FTYPE_MOSTLY_MXFP4_MOE: return GGML_TYPE_MXFP4; diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index cca08f5dc5fd..16e3c74c3d9f 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -8338,7 +8338,7 @@ static const ggml_type all_types[] = { GGML_TYPE_Q8_0, GGML_TYPE_Q1_0, GGML_TYPE_Q2_0, - GGML_TYPE_PQ2_0, + GGML_TYPE_PQ2_0, GGML_TYPE_PTQ1_0, GGML_TYPE_MXFP4, GGML_TYPE_NVFP4, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, GGML_TYPE_Q4_K, GGML_TYPE_Q5_K, @@ -8355,7 +8355,7 @@ static const ggml_type base_types[] = { GGML_TYPE_Q8_0, // for I8MM tests GGML_TYPE_Q1_0, GGML_TYPE_Q2_0, - GGML_TYPE_PQ2_0, + GGML_TYPE_PQ2_0, GGML_TYPE_PTQ1_0, GGML_TYPE_Q4_0, GGML_TYPE_Q4_1, // for I8MM tests GGML_TYPE_Q4_K, @@ -8369,7 +8369,7 @@ static const ggml_type other_types[] = { GGML_TYPE_Q8_0, GGML_TYPE_Q1_0, GGML_TYPE_Q2_0, - GGML_TYPE_PQ2_0, + GGML_TYPE_PQ2_0, GGML_TYPE_PTQ1_0, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K, GGML_TYPE_Q5_K, GGML_TYPE_Q6_K, diff --git a/tests/test-ptq1_0-cuda-dot.cpp b/tests/test-ptq1_0-cuda-dot.cpp new file mode 100644 index 000000000000..b94c279761f0 --- /dev/null +++ b/tests/test-ptq1_0-cuda-dot.cpp @@ -0,0 +1,132 @@ +// Verifies vec_dot_ptq1_0_q8_1 numerically on the host by stubbing the CUDA intrinsics. +// nvcc is unavailable here, so this cannot prove the kernel compiles, but it does prove +// the dot-product math and the element mapping agree with the CPU codec on real data. +#include +#include +#include +#include +#include + +// ---- CUDA intrinsic stubs ------------------------------------------------- +static inline int ggml_cuda_dp4a(int a, int b, int c) { // __dp4a + const int8_t* pa = (const int8_t*)&a; const int8_t* pb = (const int8_t*)&b; + return c + pa[0]*pb[0] + pa[1]*pb[1] + pa[2]*pb[2] + pa[3]*pb[3]; +} +struct half2_stub { float lo, hi; }; +static inline float __low2float(half2_stub h) { return h.lo; } + +#define QK8_1 32 +struct block_q8_1 { half2_stub ds; int8_t qs[QK8_1]; }; +static inline int get_int_b4(const int8_t* qs, int j) { // 4 bytes as an int + int v; memcpy(&v, qs + j*4, 4); return v; +} + +#define QK_PTQ1_0 128 +struct block_ptq1_0 { uint8_t qs[24]; uint8_t qh[2]; float d; }; + +// ---- transcribed from ggml-cuda/dequantize.cuh --------------------------- +static inline int ptq1_0_trit(const block_ptq1_0 * x, const int e) { + uint8_t b; int n; + if (e < 80) { b = x->qs[e & 15]; n = e >> 4; } + else if (e < 120) { const int t = e - 80; b = x->qs[16 + (t & 7)]; n = t >> 3; } + else { const int t = e - 120; b = x->qh[t & 1]; n = t >> 1; } + uint32_t v = b; + for (int i = 0; i < 4; ++i) if (i < n) v = (v * 3) & 0xFF; + return (int)((v * 3) >> 8) - 1; +} + +// ---- transcribed from ggml-cuda/vecdotq.cuh ------------------------------ +static inline float vec_dot_ptq1_0_q8_1(const void* vbq, const block_q8_1* bq8_1, + const int& kbx, const int& iqs) { + const block_ptq1_0 * bq = (const block_ptq1_0 *) vbq + kbx; + const block_q8_1 * bq8 = bq8_1 + iqs; + const int base = iqs * 32; + int sumi = 0; + for (int j = 0; j < 8; ++j) { + const int t0 = ptq1_0_trit(bq, base + j*4 + 0); + const int t1 = ptq1_0_trit(bq, base + j*4 + 1); + const int t2 = ptq1_0_trit(bq, base + j*4 + 2); + const int t3 = ptq1_0_trit(bq, base + j*4 + 3); + const int qx = (t0 & 0xFF) | ((t1 & 0xFF) << 8) | ((t2 & 0xFF) << 16) | ((t3 & 0xFF) << 24); + const int u = get_int_b4(bq8->qs, j); + sumi = ggml_cuda_dp4a(u, qx, sumi); + } + return (float) bq->d * __low2float(bq8->ds) * sumi; +} + +// ---- reference: dequantize the block, dequantize q8_1, dot in float ------ +static void ref_dequant(const block_ptq1_0* x, float* out) { + const uint8_t pow3[6]={1,3,9,27,81,243}; const size_t st[3]={32,16,8}; + int o=0; size_t j=0; + for (size_t s=0;s<3;++s){ const size_t c=st[s]; + for(; j+c<=sizeof(x->qs); j+=c) + for(size_t n=0;n<5;++n) for(size_t m=0;mqs[j+m]*pow3[n]; out[o++]=(float)((int)(((uint16_t)q*3)>>8)-1)*x->d; } + } + for(size_t n=0;n<4;++n) for(size_t h=0;h<2;++h){ + uint8_t q=x->qh[h]*pow3[n]; out[o++]=(float)((int)(((uint16_t)q*3)>>8)-1)*x->d; } +} + +int main(void) { + unsigned seed=99; + auto rnd=[&](){ seed=seed*1103515245u+12345u; return (seed>>16)&0xFFFF; }; + double worst_rel = 0.0; long checks = 0; long int_bad = 0, int_checks = 0; double worst_scaled = 0.0; + + for (int trial=0; trial<5000; ++trial) { + block_ptq1_0 w; + for (int i=0;i<24;++i) w.qs[i]=rnd()&0xFF; + for (int i=0;i<2;++i) w.qh[i]=rnd()&0xFF; + w.d = 0.01f + (rnd()%1000)/50000.0f; + + block_q8_1 y[4]; + for (int b=0;b<4;++b) { + y[b].ds.lo = 0.005f + (rnd()%1000)/80000.0f; y[b].ds.hi = 0.f; + for (int i=0;i<32;++i) y[b].qs[i]=(int8_t)((int)(rnd()%255)-127); + } + + float wf[QK_PTQ1_0]; ref_dequant(&w, wf); + + // EXACT test: the kernel's integer accumulator per chunk must equal the + // reference integer sum of trit*q8. This isolates logic from float rounding. + for (int c = 0; c < 4; ++c) { + int ref_sumi = 0; + for (int i = 0; i < 32; ++i) { + const int trit = (int) llround((double) wf[c*32+i] / (double) w.d); + ref_sumi += trit * (int) y[c].qs[i]; + } + // recompute the kernel's sumi by dividing its float result back out + const float got = vec_dot_ptq1_0_q8_1(&w, y, 0, c); + const int got_sumi = (int) llround((double) got / ((double) w.d * (double) y[c].ds.lo)); + if (ref_sumi != got_sumi) { ++int_bad; if (int_bad < 4) + printf(" INT MISMATCH trial %d chunk %d: ref %d got %d\n", trial, c, ref_sumi, got_sumi); } + ++int_checks; + } + + // reference dot over all four chunks, in float + double ref_total = 0.0; + for (int c=0;c<4;++c) + for (int i=0;i<32;++i) + ref_total += (double)wf[c*32+i] * ((double)y[c].qs[i] * (double)y[c].ds.lo); + + // kernel dot, chunk by chunk as MMVQ calls it + double got_total = 0.0; + for (int c=0;c<4;++c) got_total += vec_dot_ptq1_0_q8_1(&w, y, 0, c); + + // scale by the sum of magnitudes so cancellation in ref_total cannot inflate it + double mag = 0.0; + for (int c=0;c<4;++c) for (int i=0;i<32;++i) + mag += fabs((double)wf[c*32+i] * (double)y[c].qs[i] * (double)y[c].ds.lo); + const double denom = fabs(ref_total) > 1e-9 ? fabs(ref_total) : 1.0; + const double rel = fabs(got_total - ref_total) / denom; + if (rel > worst_rel) worst_rel = rel; + if (mag > 0) { const double sc = fabs(got_total - ref_total)/mag; if (sc > worst_scaled) worst_scaled = sc; } + ++checks; + } + printf(" dot products compared : %ld (4 chunks each)\n", checks); + printf(" worst relative error : %.3e\n", worst_rel); + printf(" worst err / sum|terms| : %.3e (immune to cancellation)\n", worst_scaled); + printf(" exact integer checks : %ld, mismatches %ld\n", int_checks, int_bad); + if (int_bad == 0) printf(" LOGIC EXACT: integer accumulator matches reference on every chunk\n"); + else printf(" LOGIC BUG in the kernel\n"); + return int_bad != 0; +} diff --git a/tests/test-ptq1_0-element-map.cpp b/tests/test-ptq1_0-element-map.cpp new file mode 100644 index 000000000000..d8cf77e10549 --- /dev/null +++ b/tests/test-ptq1_0-element-map.cpp @@ -0,0 +1,64 @@ +// Verifies the CUDA device accessor's element mapping matches the CPU codec exactly. +// The CUDA kernels cannot be compiled here (no nvcc), but this is where a mapping bug +// would hide, and it is pure integer logic, so it is checkable on the host. +#include +#include +#include +#include + +#define QK_PTQ1_0 128 +struct block_ptq1_0 { uint8_t qs[24]; uint8_t qh[2]; uint16_t d; }; + +// --- transcribed verbatim from ggml-cuda/dequantize.cuh --- +static int ptq1_0_trit(const block_ptq1_0 * x, const int e) { + uint8_t b; int n; + if (e < 80) { b = x->qs[e & 15]; n = e >> 4; } + else if (e < 120) { const int t = e - 80; b = x->qs[16 + (t & 7)]; n = t >> 3; } + else { const int t = e - 120; b = x->qh[t & 1]; n = t >> 1; } + uint32_t v = b; + for (int i = 0; i < 4; ++i) if (i < n) v = (v * 3) & 0xFF; + return (int)((v * 3) >> 8) - 1; +} + +// --- the CPU reference traversal from ggml-quants.c dequantize_row_ptq1_0 --- +static void cpu_ref(const block_ptq1_0 * x, int * out) { + const uint8_t pow3[6] = {1,3,9,27,81,243}; + const size_t stages[3] = {32,16,8}; + int o = 0; size_t j = 0; + for (size_t s = 0; s < 3; ++s) { + const size_t c = stages[s]; + for (; j + c <= sizeof(x->qs); j += c) + for (size_t n = 0; n < 5; ++n) + for (size_t m = 0; m < c; ++m) { + uint8_t q = x->qs[j+m] * pow3[n]; + out[o++] = (int)(((uint16_t)q * 3) >> 8) - 1; + } + } + for (size_t n = 0; n < 4; ++n) + for (size_t h = 0; h < sizeof(x->qh); ++h) { + uint8_t q = x->qh[h] * pow3[n]; + out[o++] = (int)(((uint16_t)q * 3) >> 8) - 1; + } +} + +int main(void) { + unsigned seed = 12345; + long bad = 0, total = 0; + for (int trial = 0; trial < 20000; ++trial) { + block_ptq1_0 blk; + for (int i = 0; i < 24; ++i) { seed = seed*1103515245u+12345u; blk.qs[i] = (seed>>16)&0xFF; } + for (int i = 0; i < 2; ++i) { seed = seed*1103515245u+12345u; blk.qh[i] = (seed>>16)&0xFF; } + int ref[QK_PTQ1_0]; cpu_ref(&blk, ref); + for (int e = 0; e < QK_PTQ1_0; ++e) { + ++total; + if (ptq1_0_trit(&blk, e) != ref[e]) { + if (bad < 5) printf(" MISMATCH blk%d e=%d cuda=%d cpu=%d\n", trial, e, ptq1_0_trit(&blk,e), ref[e]); + ++bad; + } + } + } + printf(" checked %ld element positions across 20000 random blocks\n", total); + printf(" mismatches: %ld\n", bad); + printf(" %s\n", bad==0 ? "CUDA element mapping MATCHES the CPU codec exactly" : "MAPPING BUG"); + return bad != 0; +} diff --git a/tests/test-quantize-fns.cpp b/tests/test-quantize-fns.cpp index 2811507a3cd2..4f77fb9701c7 100644 --- a/tests/test-quantize-fns.cpp +++ b/tests/test-quantize-fns.cpp @@ -160,6 +160,7 @@ static int test_vec_dot_q(bool verbose) { type == GGML_TYPE_TQ2_0 ? MAX_QUANTIZATION_TOTAL_ERROR_TERNARY : type == GGML_TYPE_Q2_0 ? MAX_QUANTIZATION_TOTAL_ERROR_TERNARY : type == GGML_TYPE_PQ2_0 ? MAX_QUANTIZATION_TOTAL_ERROR_TERNARY : + type == GGML_TYPE_PTQ1_0 ? MAX_QUANTIZATION_TOTAL_ERROR_TERNARY : type == GGML_TYPE_Q2_K ? MAX_QUANTIZATION_TOTAL_ERROR_2BITS : type == GGML_TYPE_IQ2_S ? MAX_QUANTIZATION_TOTAL_ERROR_2BITS : type == GGML_TYPE_Q3_K ? MAX_QUANTIZATION_TOTAL_ERROR_3BITS : @@ -185,7 +186,7 @@ static int test_vec_dot_q(bool verbose) { ? MAX_DOT_PRODUCT_ERROR_LOWBIT : type == GGML_TYPE_Q1_0 ? MAX_DOT_PRODUCT_ERROR_BINARY - : type == GGML_TYPE_TQ1_0 || type == GGML_TYPE_TQ2_0 || type == GGML_TYPE_Q2_0 || type == GGML_TYPE_PQ2_0 + : type == GGML_TYPE_TQ1_0 || type == GGML_TYPE_TQ2_0 || type == GGML_TYPE_Q2_0 || type == GGML_TYPE_PQ2_0 || type == GGML_TYPE_PTQ1_0 ? MAX_DOT_PRODUCT_ERROR_TERNARY : type == GGML_TYPE_NVFP4 ? MAX_DOT_PRODUCT_ERROR_FP4 diff --git a/tools/quantize/quantize.cpp b/tools/quantize/quantize.cpp index ff54010f247a..fe668c3ee123 100644 --- a/tools/quantize/quantize.cpp +++ b/tools/quantize/quantize.cpp @@ -35,6 +35,7 @@ static const std::vector QUANT_OPTIONS = { { "Q1_0", LLAMA_FTYPE_MOSTLY_Q1_0, " 1.125 bpw quantization", }, { "Q2_0", LLAMA_FTYPE_MOSTLY_Q2_0, " 2.25 bpw quantization (group 64)", }, { "PQ2_0", LLAMA_FTYPE_MOSTLY_PQ2_0, " 2.13 bpw quantization (group 128, Prism)", }, + { "PTQ1_0", LLAMA_FTYPE_MOSTLY_PTQ1_0, " 1.75 bpw ternarization (group 128, Prism)", }, { "Q4_0", LLAMA_FTYPE_MOSTLY_Q4_0, " 4.34G, +0.4685 ppl @ Llama-3-8B", }, { "Q4_1", LLAMA_FTYPE_MOSTLY_Q4_1, " 4.78G, +0.4511 ppl @ Llama-3-8B", }, { "MXFP4_MOE",LLAMA_FTYPE_MOSTLY_MXFP4_MOE," MXFP4 MoE", },