From e19819227256ede294a825bbbc94a0be54dcf54f Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Wed, 2 Sep 2026 20:34:52 -0700 Subject: [PATCH 01/14] ggml: add PTQ1_0, Prism ternary at group 128 (1.75 bpw) Upstream TQ1_0 packs ternary weights base-3 at 5 trits per byte, but carries a single fp16 scale per 256 weights. The ternary checkpoints this targets carry a scale per 128, so a 256-wide scale has to discard one of the two group scales it straddles. Measured on a ternary checkpoint of this family: every trit survives and the zero pattern is preserved exactly, but the smaller-alpha group of each block is reconstructed at the larger alpha, giving 6.43% weight error that passes into activations at full strength because it is multiplicative on the weight rather than noise. PTQ1_0 keeps TQ1_0's trit packing and our group of 128: qs[24] + qh[2] + fp16 d = 28 bytes / 128 weights = 1.7500 bpw against PQ2_0's 34 bytes = 2.1250 bpw, so 17.6% smaller with no loss of information: PQ2_0 spends two bits on a value carrying log2(3) = 1.585 bits. test-quantize-fns reports ptq1_0 identical to pq2_0 on both metrics, absolute quantization error 0.008678 and dot product error 0.141111, versus tq1_0's 0.008681 and 0.141345. A round trip of real folded ternary weights is bit-identical to the PQ2_0 reconstruction, so the only error either format carries is the shared fp16 scale rounding. A distinct type id (143) rather than changing TQ1_0 in place, for the same reason PQ2_0 is distinct from upstream's group-64 Q2_0: TQ1_0 is a wire format, and reinterpreting its block width would silently misread every existing TQ1_0 gguf. TQ1_0's qs staging is fixed at 32-then-16 bytes, which cannot cover 24, so the stages are generalised to 32/16/8. At TQ1_0's 48-byte qs that reduces to exactly its original behaviour. Scope: codec, ggml traits, CPU vec_dot against Q8_0, and the ftype plumbing so llama-quantize can target it. The vec_dot decodes a block to element order before the dot, sharing the traversal with dequantize so the two cannot drift; it is generic C on every arch, with no SIMD path yet. CUDA, Metal and Vulkan kernels, and the repack gemv/gemm paths, are not implemented. --- ggml/include/ggml.h | 4 +- ggml/src/ggml-common.h | 13 ++++ ggml/src/ggml-cpu/arch-fallback.h | 15 +++++ ggml/src/ggml-cpu/ggml-cpu.c | 6 ++ ggml/src/ggml-cpu/quants.c | 72 +++++++++++++++++++++ ggml/src/ggml-cpu/quants.h | 3 + ggml/src/ggml-quants.c | 101 ++++++++++++++++++++++++++++++ ggml/src/ggml-quants.h | 3 + ggml/src/ggml.c | 8 +++ include/llama.h | 1 + src/llama-model-loader.cpp | 2 + src/llama-quant.cpp | 4 +- tests/test-quantize-fns.cpp | 3 +- tools/quantize/quantize.cpp | 1 + 14 files changed, 233 insertions(+), 3 deletions(-) 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..0f0fe6b596c5 100644 --- a/ggml/src/ggml-common.h +++ b/ggml/src/ggml-common.h @@ -204,6 +204,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/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-quants.c b/ggml/src/ggml-quants.c index a5f8f9c352f3..d53ba2a7b998 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); 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.c b/ggml/src/ggml.c index 6849721d8c60..e06e07312931 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, 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-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", }, From dce9a4b90d8ab02d95a47fcdea8175b05066ecfe Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Wed, 2 Sep 2026 20:50:34 -0700 Subject: [PATCH 02/14] ggml: wire PTQ1_0 into quantize dispatch, ftype map and row validation Three plumbing gaps the traits entry alone did not cover, each of which failed end-to-end quantization in a different way: ggml_quantize_chunk - without a dispatch case the returned byte count did not match the row size and GGML_ASSERT(result == nrows * row_size) aborted a worker thread mid-tensor, leaving a truncated output file behind. ggml_ftype -> wtype - the ftype could not resolve to a ggml type. validate_row_data - "invalid type 143", so quantization refused the result it had just produced. Verified end to end on a 1.7B ternary model with the embedding and output tensors pinned to f16 in both runs, so only the quantized tensors differ: quantized tensors PQ2_0 357.47 MiB -> PTQ1_0 294.47 MiB, 17.62% smaller (1.75/2.125 bpw predicts 17.65%) perplexity, 4 chunks at ctx 512, identical to every printed digit: per chunk 11.3796 / 14.7182 / 14.5014 / 12.6944 for both, final 12.6944 +/- 1.12596 for both The perplexity run goes through the new CPU vec_dot, so that path is exercised rather than only the codec. --- ggml/src/ggml-quants.c | 4 ++++ ggml/src/ggml.c | 2 ++ 2 files changed, 6 insertions(+) diff --git a/ggml/src/ggml-quants.c b/ggml/src/ggml-quants.c index d53ba2a7b998..bfb3892a599c 100644 --- a/ggml/src/ggml-quants.c +++ b/ggml/src/ggml-quants.c @@ -5712,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.c b/ggml/src/ggml.c index e06e07312931..2e7416070d04 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -1451,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; @@ -8058,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; From e5bb8d3a0c1869a1905f24017e81154ba2a72a46 Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Wed, 2 Sep 2026 21:00:55 -0700 Subject: [PATCH 03/14] =?UTF-8?q?metal:=20PTQ1=5F0=20kernels=20=E2=80=94?= =?UTF-8?q?=20correct,=20but=204.2x=20slower=20decode=20than=20PQ2=5F0?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Adds the Metal path for PTQ1_0: an element accessor for the base-3 layout, both dequantize entry points, a mat-vec kernel, mul_mm and mul_mm_id instantiations for f32 and f16, get_rows, the simdgroup constants and both pipeline-selection switches. Also enumerates PTQ1_0 in test-backend-ops, which otherwise silently skipped the new type and reported a clean pass. Correct: 1251/1251 MUL_MAT tests pass on Apple M5 Pro. End to end on a 1.7B ternary model the Metal perplexity is identical to PQ2_0 to every printed digit, 12.6935 +/- 1.12567 for both. Slow, and the reason is structural. Measured on M5 Pro: PQ2_0 PTQ1_0 pp128 4505 tok/s 2590 tok/s 0.57x tg32 182.5 tok/s 43.6 tok/s 0.24x Base-3 packing makes element order non-positional, so a thread owning 16 contiguous weights cannot index bytes directly, and threads owning different trit indices reload the same bytes. Hoisting the trit factor out of the inner loop for spans below element 80 took decode from 29.8 to 43.6 tok/s, but the remaining gap needs the thread mapping changed so one thread consumes all five trits of a byte it loads once, instead of five threads each loading it for one trit. This reproduces what the earlier TQ1_0 work already found on other backends: base-3 is ALU-bound rather than bandwidth-bound, and the packing win does not pay for itself at decode without that restructure. The 17.6% file-size reduction is real and lossless; the speed is not there yet. --- ggml/src/ggml-metal/ggml-metal-device.cpp | 10 ++ ggml/src/ggml-metal/ggml-metal-impl.h | 2 + ggml/src/ggml-metal/kernels/dequantize.h | 44 +++++++++ ggml/src/ggml-metal/kernels/mul_mm.metal | 4 + ggml/src/ggml-metal/kernels/mul_mv.metal | 105 +++++++++++++++++++++ ggml/src/ggml-metal/kernels/quantize.metal | 1 + tests/test-backend-ops.cpp | 6 +- 7 files changed, 169 insertions(+), 3 deletions(-) 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-impl.h b/ggml/src/ggml-metal/ggml-metal-impl.h index 836cbbcb5c9e..2ca1446c1b6b 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 8 +#define N_SG_PTQ1_0 2 #define N_R0_Q4_0 4 #define N_SG_Q4_0 2 diff --git a/ggml/src/ggml-metal/kernels/dequantize.h b/ggml/src/ggml-metal/kernels/dequantize.h index 1254d190c9b1..fbb0f182ea4d 100644 --- a/ggml/src/ggml-metal/kernels/dequantize.h +++ b/ggml/src/ggml-metal/kernels/dequantize.h @@ -161,6 +161,50 @@ void dequantize_pq2_0_t4(device const block_pq2_0 * xb, short il, thread type4 & reg = (type4) reg_f; } +// 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) { + const uchar pow3[5] = {1, 3, 9, 27, 81}; + uchar q; + if (e < 80) { // qs[0..15], chunk of 16, 5 trits per byte + q = xb->qs[e & 15] * pow3[e >> 4]; + } else if (e < 120) { // qs[16..23], chunk of 8 + const int t = e - 80; + q = xb->qs[16 + (t & 7)] * pow3[t >> 3]; + } else { // qh[0..1], 4 trits per byte + const int t = e - 120; + q = xb->qh[t & 1] * pow3[t >> 1]; + } + return (float) ((int) (((ushort) q * 3) >> 8) - 1); +} + +template +void dequantize_ptq1_0(device const block_ptq1_0 * xb, short il, thread type4x4 & reg) { + const float d = xb->d; + const int base = il * 16; + + float4x4 reg_f; + 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..cc09e0a1ed8e 100644 --- a/ggml/src/ggml-metal/kernels/mul_mv.metal +++ b/ggml/src/ggml-metal/kernels/mul_mv.metal @@ -510,6 +510,108 @@ kernel void kernel_mul_mv_q2_0_f32( kernel_mul_mv_q2_0_f32_impl(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg); } +// A 16-element span starting at a multiple of 16 below element 80 lies entirely +// inside one trit index of the 16-byte chunk, so the branch and the pow3 factor +// hoist out of the loop and the 16 byte loads are the only per-element work. +// The spans at 80 and above straddle trit indices, so they take the general path. +inline float ptq1_0_dot_y(device const block_ptq1_0 * qb, thread float * yl, int il) { + const uchar pow3[5] = {1, 3, 9, 27, 81}; + + float sum = 0.f; + + if (il < 80) { + const uchar c = pow3[il >> 4]; + device const uchar * qs = qb->qs; + FOR_UNROLL (short i = 0; i < 16; ++i) { + const uchar q = qs[i] * c; + sum += (float) ((int) (((ushort) q * 3) >> 8) - 1) * yl[i]; + } + } else { + for (short i = 0; i < 16; ++i) { + sum += ptq1_0_elem(qb, il + i) * yl[i]; + } + } + + return sum * (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}; + + // group 128: 8 sub-blocks of 16 weights per PTQ1_0 block + const short ix = (tiisg/8); + const short il = (tiisg%8)*16; + + device const float * yb = y + ix*QK_PTQ1_0 + il; + + for (int ib = ix; ib < nb; ib += N_SIMDWIDTH/8) { + FOR_UNROLL (short i = 0; i < 16; i++) { + yl[i] = yb[i]; + } + + FOR_UNROLL (short row = 0; row < nr0; row++) { + sumf[row] += ptq1_0_dot_y(ax[row] + ib, yl, il); + } + + 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 +1097,9 @@ 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_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/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, From 745bf55fc8c4c048194d41e4cce2797fdac269d1 Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Wed, 2 Sep 2026 21:14:19 -0700 Subject: [PATCH 04/14] metal: PTQ1_0 decode 29.8 -> 165.3 tok/s, now 0.89x of PQ2_0 Four changes, each measured on Apple M5 Pro against a 1.7B ternary model, with MUL_MAT staying 1251/1251 and Metal perplexity identical to PQ2_0 throughout: 1. byte-owning threads 29.8 -> 136.4 a thread now consumes all five trits of a byte it loads once, instead of eight threads each reloading the same qs bytes for one trit index; the block's 26 bytes are read once total 2. y staged in registers 136.4 -> 143.5 y was being re-read from device memory once per row; it is now gathered once per block and reused, with the ternary offset folded into a single sumy subtraction 3. N_R0 8 -> 4 143.5 -> 163.8 rows-per-thread was the real limiter: 16 rows gives 94 tok/s, 8 gives 144, 4 gives 164, 2 gives 163 4. balanced qh 163.8 -> 165.3 qh's 8 elements went to threads 0 and 1 only, so those did 19 elements while the rest did 15 and the whole simdgroup waited; one element per thread makes it 16 each and drops yl to 16 Trit extraction is the base-3 remainder recurrence, t = (v*3)>>8 with v = (v*3)&0xFF, two integer ops per trit with no table. A 256-entry lookup was tried and is not better for decode: it adds a dependent load per byte, 26 extra per block against PQ2_0's 32 total loads. It does help the mul_mm path, which must produce 16 contiguous elements and cannot use the byte-owning mapping, so dequantize keeps the table. Final, PQ2_0 vs PTQ1_0 on M5 Pro: pp128 4541 -> 3998 tok/s 0.88x tg32 186 -> 165 tok/s 0.89x size 357 -> 294 MiB 17.6% smaller, and lossless Prefill's remaining gap is structural: mul_mm's dequantize contract is 16 contiguous elements, so five threads still touch the same 16 bytes for different trit indices. Closing it needs mul_mm itself to stage a decoded block in threadgroup memory, which is a larger change than this. --- ggml/src/ggml-metal/ggml-metal-impl.h | 2 +- ggml/src/ggml-metal/kernels/dequantize.h | 59 ++++++++++--- ggml/src/ggml-metal/kernels/mul_mv.metal | 100 +++++++++++++++++------ 3 files changed, 126 insertions(+), 35 deletions(-) diff --git a/ggml/src/ggml-metal/ggml-metal-impl.h b/ggml/src/ggml-metal/ggml-metal-impl.h index 2ca1446c1b6b..d1e9aa8fd8e9 100644 --- a/ggml/src/ggml-metal/ggml-metal-impl.h +++ b/ggml/src/ggml-metal/ggml-metal-impl.h @@ -29,7 +29,7 @@ #define N_R0_PQ2_0 8 #define N_SG_PQ2_0 2 -#define N_R0_PTQ1_0 8 +#define N_R0_PTQ1_0 4 #define N_SG_PTQ1_0 2 #define N_R0_Q4_0 4 diff --git a/ggml/src/ggml-metal/kernels/dequantize.h b/ggml/src/ggml-metal/kernels/dequantize.h index fbb0f182ea4d..1e01dcd81d9e 100644 --- a/ggml/src/ggml-metal/kernels/dequantize.h +++ b/ggml/src/ggml-metal/kernels/dequantize.h @@ -161,35 +161,74 @@ 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) { - const uchar pow3[5] = {1, 3, 9, 27, 81}; - uchar q; + uchar b; + short n; if (e < 80) { // qs[0..15], chunk of 16, 5 trits per byte - q = xb->qs[e & 15] * pow3[e >> 4]; + b = xb->qs[e & 15]; n = e >> 4; } else if (e < 120) { // qs[16..23], chunk of 8 const int t = e - 80; - q = xb->qs[16 + (t & 7)] * pow3[t >> 3]; + b = xb->qs[16 + (t & 7)]; n = t >> 3; } else { // qh[0..1], 4 trits per byte const int t = e - 120; - q = xb->qh[t & 1] * pow3[t >> 1]; + b = xb->qh[t & 1]; n = t >> 1; } - return (float) ((int) (((ushort) q * 3) >> 8) - 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; - const int base = il * 16; float4x4 reg_f; - 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; + + // 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; } diff --git a/ggml/src/ggml-metal/kernels/mul_mv.metal b/ggml/src/ggml-metal/kernels/mul_mv.metal index cc09e0a1ed8e..c87679038990 100644 --- a/ggml/src/ggml-metal/kernels/mul_mv.metal +++ b/ggml/src/ggml-metal/kernels/mul_mv.metal @@ -510,29 +510,61 @@ kernel void kernel_mul_mv_q2_0_f32( kernel_mul_mv_q2_0_f32_impl(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg); } -// A 16-element span starting at a multiple of 16 below element 80 lies entirely -// inside one trit index of the 16-byte chunk, so the branch and the pow3 factor -// hoist out of the loop and the 16 byte loads are the only per-element work. -// The spans at 80 and above straddle trit indices, so they take the general path. -inline float ptq1_0_dot_y(device const block_ptq1_0 * qb, thread float * yl, int il) { - const uchar pow3[5] = {1, 3, 9, 27, 81}; - - float sum = 0.f; +// 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; + } + } - if (il < 80) { - const uchar c = pow3[il >> 4]; - device const uchar * qs = qb->qs; - FOR_UNROLL (short i = 0; i < 16; ++i) { - const uchar q = qs[i] * c; - sum += (float) ((int) (((ushort) q * 3) >> 8) - 1) * yl[i]; + { + 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; } - } else { - for (short i = 0; i < 16; ++i) { - sum += ptq1_0_elem(qb, il + i) * yl[i]; + } + + // 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 sum * (float) qb->d; + return (acc - sumy) * (float) qb->d; } template @@ -571,19 +603,39 @@ void kernel_mul_mv_ptq1_0_f32_impl( float yl[16]; float sumf[nr0] = {0.f}; - // group 128: 8 sub-blocks of 16 weights per PTQ1_0 block + // 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 il = (tiisg%8)*16; + const short it = (tiisg%8); - device const float * yb = y + ix*QK_PTQ1_0 + il; + device const float * yb = y + ix*QK_PTQ1_0; for (int ib = ix; ib < nb; ib += N_SIMDWIDTH/8) { - FOR_UNROLL (short i = 0; i < 16; i++) { - yl[i] = yb[i]; + // 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_y(ax[row] + ib, yl, il); + sumf[row] += ptq1_0_dot_reg(ax[row] + ib, yl, sumy, it); } yb += QK_PTQ1_0 * (N_SIMDWIDTH/8); From e2ecfa6150c3ca74a8b0b3aecfff456453d5bd62 Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Wed, 2 Sep 2026 21:17:58 -0700 Subject: [PATCH 05/14] =?UTF-8?q?cuda:=20PTQ1=5F0=20dequantize=20and=20MMV?= =?UTF-8?q?Q=20paths=20=E2=80=94=20NOT=20COMPILED,=20no=20toolchain=20here?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Adds the CUDA surface for PTQ1_0, mirroring PQ2_0: type traits, a device trit accessor, the float2 dequantize entry point, the six dequant dispatch sites in convert.cu, get_rows, the MMVQ dot with its VDR constant and three dispatch sites, and both supports_op type lists. **This code has never been compiled or executed.** The machine it was written on is an Apple M5 Pro with no CUDA toolchain, so nvcc was unavailable. It must be built and run on a CUDA host before anyone relies on it. What *is* verified is the part where a bug would actually hide. The device accessor's element mapping is pure integer logic, so it was transcribed verbatim to the host and checked against the CPU codec's own traversal over 2,560,000 element positions across 20,000 random blocks: zero mismatches. That check is kept as tests/test-ptq1_0-element-map.cpp so it can be rerun. Trit extraction uses the same base-3 remainder recurrence as the Metal path, t = (v*3)>>8 with v = (v*3)&0xFF. The MMVQ dot gathers four trits at a time and packs them into an int as signed int8 lanes so accumulation still goes through dp4a, since base-3 packing leaves no byte-aligned run inside a 32-element chunk. MMQ is deliberately not implemented. Types without it fall back to dequantize plus cuBLAS for prefill, which is correct if not optimal, and writing an MMQ tile loader blind against tile layouts I cannot test would be worse than leaving the fallback in place. --- ggml/src/ggml-common.h | 2 + ggml/src/ggml-cuda/common.cuh | 7 ++++ ggml/src/ggml-cuda/convert.cu | 12 ++++++ ggml/src/ggml-cuda/dequantize.cuh | 36 +++++++++++++++++ ggml/src/ggml-cuda/getrows.cu | 3 ++ ggml/src/ggml-cuda/ggml-cuda.cu | 2 + ggml/src/ggml-cuda/mmvq.cu | 8 ++++ ggml/src/ggml-cuda/vecdotq.cuh | 33 ++++++++++++++++ tests/test-ptq1_0-element-map.cpp | 64 +++++++++++++++++++++++++++++++ 9 files changed, 167 insertions(+) create mode 100644 tests/test-ptq1_0-element-map.cpp diff --git a/ggml/src/ggml-common.h b/ggml/src/ggml-common.h index 0f0fe6b596c5..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)) diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index b2d82beaf7c8..652ba8d00339 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -992,6 +992,13 @@ struct ggml_cuda_type_traits { static constexpr int qi = QI_PQ2_0; }; +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..c8c21ee8277f 100644 --- a/ggml/src/ggml-cuda/dequantize.cuh +++ b/ggml/src/ggml-cuda/dequantize.cuh @@ -43,6 +43,42 @@ static __device__ __forceinline__ void dequantize_q2_0(const void * vx, const in v.y = (c1 - 1) * d; } +// 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; +} + +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..c3f365458df2 100644 --- a/ggml/src/ggml-cuda/getrows.cu +++ b/ggml/src/ggml-cuda/getrows.cu @@ -323,6 +323,9 @@ static void ggml_cuda_get_rows_switch_src0_type( case GGML_TYPE_PQ2_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); + 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, 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..5ac3eaf30c4d 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; @@ -1120,6 +1122,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..6484b15506ba 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 1 // one 32-element chunk at a time (4 chunks per 128 block) #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,6 +803,38 @@ static __device__ __forceinline__ float vec_dot_q2_0_q8_1( return d2 * d8 * sumi; } +// PTQ1_0 x Q8_1. iqs selects one of the four 32-element chunks of the 128-weight +// block, pairing with one Q8_1 block, same convention as PQ2_0. Trits are gathered +// individually because base-3 packing gives no byte-aligned run inside a chunk, then +// packed four at a time into an int so the accumulation can still use dp4a. +// NOT COMPILED OR RUN: no CUDA toolchain was available where this was written. +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) { + + 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; +#pragma unroll + 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); + + // four signed int8 lanes, -1/0/1 + 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); + } + + const float d8 = __low2float(bq8->ds); + return (float) bq->d * 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) { 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; +} From 9208cdc3903791354910499d2b6572e31cf7f5c7 Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Wed, 2 Sep 2026 21:30:16 -0700 Subject: [PATCH 06/14] metal: tune PTQ1_0 to N_SG=1, and record what optimising it ruled out N_R0=4, N_SG=1. Decode lands at 162-166 tok/s across runs against PQ2_0's 186, so 0.87-0.89x; the spread between N_SG 1 and 2 is inside run-to-run noise (+/- 3.4 on a single measurement). Five further optimisations were tried and MEASURED, and four of them lost. Kept here so nobody spends the same day twice: 256-entry trit lookup table 142.1 vs 143.5 recurrence no decode gain (helps mul_mm, which keeps it) independent pow3 extraction 159.8 vs 166.0 recurrence WORSE contiguous byte ownership 162.0 vs 165.3 scattered WORSE N_R0 = 3 / 5 / 6 129.8 / 156.7 / 149.6 all worse than 4 N_SG = 8 161.1 worse than 1-2 What that rules out. The lookup table shows decode is not bound by trit arithmetic. The pow3 variant shows it is not bound by the recurrence's five-deep serial dependency either -- four rows per thread already supply enough independent work, so fewer ops beat more ILP. Contiguous byte runs made it worse because thread 5 straddles the qs 16/8 chunk boundary, and the resulting per-byte branch cost more than consolidating the loads saved. Where the remaining 11% actually lives, per 128 weights: PQ2_0 128 x (shift, mask, cvt, fma) = 512 ops, 34 B read PTQ1_0 128 x (mul, shift, cvt, fma, mask-for-next) = 640 ops, 28 B read ALU 1.25x, bytes 0.82x A purely ALU-bound kernel would sit at 0.80x and a purely bandwidth-bound one at 1.21x. Measured 0.89x sits between them, which is about where this instruction and byte mix should land. Extracting a base-3 digit is inherently a multiply where a 2-bit field is a shift, so this gap is the price of packing at the 1.585-bit information limit rather than on a 2-bit boundary, not an implementation defect. The one lever left with a real mechanism behind it is prefill, not decode. In decode each row has its own weights so there is nothing to share, but mul_mm reuses a weight tile across many output columns, and its dequantize contract forces 16 contiguous elements, so five threads still touch the same 16 bytes. Staging a cooperatively decoded block in threadgroup memory would remove that redundancy outright. Prefill is at 0.87x (3946 vs 4492). Perplexity remains identical to PQ2_0 at 12.6935 +/- 1.12567, MUL_MAT 1251/1251. --- ggml/src/ggml-metal/ggml-metal-impl.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-metal/ggml-metal-impl.h b/ggml/src/ggml-metal/ggml-metal-impl.h index d1e9aa8fd8e9..1151e5099fc8 100644 --- a/ggml/src/ggml-metal/ggml-metal-impl.h +++ b/ggml/src/ggml-metal/ggml-metal-impl.h @@ -30,7 +30,7 @@ #define N_R0_PQ2_0 8 #define N_SG_PQ2_0 2 #define N_R0_PTQ1_0 4 -#define N_SG_PTQ1_0 2 +#define N_SG_PTQ1_0 1 #define N_R0_Q4_0 4 #define N_SG_Q4_0 2 From f844a75709568de6bcd76705f0f06b0a1b89d990 Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Wed, 2 Sep 2026 21:41:50 -0700 Subject: [PATCH 07/14] metal, cpu: close PTQ1_0 type-list gaps, add the missing mul_mv_ext r1_5 Self-review compared every switch that names PQ2_0 against those naming PTQ1_0 and found the new type absent from ten support enumerations: seven in ggml-cpu/ops.cpp, including ggml_compute_forward_get_rows_f32, one in ggml-metal-ops.cpp and two in ggml-metal-device.m. PTQ1_0 is traits-compatible with PQ2_0 there -- same block size, same to_float and from_float shape -- so it belongs in each. Closing those gaps then exposed a real defect. Adding the type to the Metal support lists enabled the mul_mv_ext path for wider n, which immediately failed with "kernel not found in any metal library" for kernel_mul_mv_ext_ptq1_0_f32_r1_5: PQ2_0 instantiates r1_2 through r1_5 and only r1_2 through r1_4 had been written, because the grep that enumerated them was truncated by head. Instantiating r1_5 fixes it, and the instantiation set is now derived from PQ2_0's rather than transcribed. Worth noting the earlier 1251/1251 passes were not wrong, they were narrow: with the type missing from the support lists the wider-n path was never selected, so the absent kernel could not be reached. The gap and the defect concealed each other. After: MUL_MAT 1251/1251 on M5 Pro, Backend MTL0 OK, Metal perplexity identical to PQ2_0 at 12.6935 +/- 1.12567, decode 165 tok/s and prefill 3947. --- ggml/src/ggml-cpu/ops.cpp | 7 +++++++ ggml/src/ggml-metal/ggml-metal-device.m | 2 ++ ggml/src/ggml-metal/ggml-metal-ops.cpp | 1 + ggml/src/ggml-metal/kernels/mul_mv.metal | 1 + 4 files changed, 11 insertions(+) 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-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-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/mul_mv.metal b/ggml/src/ggml-metal/kernels/mul_mv.metal index c87679038990..222902cc7e0b 100644 --- a/ggml/src/ggml-metal/kernels/mul_mv.metal +++ b/ggml/src/ggml-metal/kernels/mul_mv.metal @@ -1152,6 +1152,7 @@ template [[host_name("kernel_mul_mv_ext_pq2_0_f32_r1_4")]] kernel mul_mv_ext_q4 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>; From 63857492b4d22ef3691450c49ef035f18ca90733 Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Wed, 2 Sep 2026 22:07:39 -0700 Subject: [PATCH 08/14] tests: verify the CUDA MMVQ dot numerically on the host The CUDA path still cannot be compiled here, but its arithmetic can be checked. This stubs the three intrinsics the kernel uses -- __dp4a, __low2float and the 4-byte int load -- transcribes vec_dot_ptq1_0_q8_1 and the device trit accessor verbatim, and compares them against a float reference built by dequantising the block and the q8_1 activations independently. Result over 5000 random weight blocks with random q8_1 activations: exact integer-accumulator checks 20000, mismatches 0 error / sum of term magnitudes 2.399e-08 The integer test is the one that matters: the kernel accumulates in int via dp4a before scaling, so if the trit extraction, the int8 lane packing or the chunk-to- q8_1 pairing were wrong, that accumulator would differ. It does not, on any of the 20000 chunks. A first pass reported 7.062e-05 "relative error" and looked like a failure. That was cancellation: the reference total is a signed sum that can land near zero, so dividing by it inflates the ratio. Measured against the sum of term magnitudes instead, which cancellation cannot distort, the error is 2.4e-08, i.e. float32 rounding of the four per-chunk products. Together with tests/test-ptq1_0-element-map.cpp this leaves only compilation unproven for the CUDA path, not correctness of its math. --- tests/test-ptq1_0-cuda-dot.cpp | 132 +++++++++++++++++++++++++++++++++ 1 file changed, 132 insertions(+) create mode 100644 tests/test-ptq1_0-cuda-dot.cpp 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; +} From 10462d037e3f27f63ffc0dc7670a5c019d1ee131 Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Wed, 2 Sep 2026 23:19:53 -0700 Subject: [PATCH 09/14] =?UTF-8?q?cuda:=20move=20ptq1=5F0=5Ftrit=20to=20com?= =?UTF-8?q?mon.cuh=20so=20vecdotq.cuh=20can=20see=20it=20=E2=80=94=20BUILD?= =?UTF-8?q?S=20NOW?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit First compile of the CUDA path, on an RTX 5090 with CUDA 12.8 and CMAKE_CUDA_ARCHITECTURES=120, failed with vecdotq.cuh(822): error: identifier "ptq1_0_trit" is undefined The accessor was defined in dequantize.cuh, which vecdotq.cuh does not include. PQ2_0's dot never needed a shared helper -- it decodes 2-bit fields inline with byte_perm -- so nothing in the file it was mirrored from exposed the dependency. Both headers include common.cuh, so the helper belongs there. Verified on the 5090 after the fix: test-backend-ops -o MUL_MAT 1283/1283 passed, Backend CUDA0: OK 45 ptq1_0 cases OK, 0 failures host element-map test 2,560,000 positions, 0 mismatches host MMVQ dot test 20,000 integer checks, 0 mismatches The two host tests were already green before this, which is the point worth recording: they verified the arithmetic and could not have caught a missing declaration. Compilation was the only gate that would find this, and the earlier commit saying the CUDA path had never been built was accurate about the risk. --- ggml/src/ggml-cuda/common.cuh | 27 +++++++++++++++++++++++++++ ggml/src/ggml-cuda/dequantize.cuh | 26 -------------------------- 2 files changed, 27 insertions(+), 26 deletions(-) diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 652ba8d00339..ef929d3d7842 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -992,6 +992,33 @@ 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; diff --git a/ggml/src/ggml-cuda/dequantize.cuh b/ggml/src/ggml-cuda/dequantize.cuh index c8c21ee8277f..7a1deec8b829 100644 --- a/ggml/src/ggml-cuda/dequantize.cuh +++ b/ggml/src/ggml-cuda/dequantize.cuh @@ -43,32 +43,6 @@ static __device__ __forceinline__ void dequantize_q2_0(const void * vx, const in v.y = (c1 - 1) * d; } -// 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; -} 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; From 4d583a9f3bed1d6185689ca59f20607d473d5ced Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Wed, 2 Sep 2026 23:46:15 -0700 Subject: [PATCH 10/14] =?UTF-8?q?vulkan:=20PTQ1=5F0=20support,=20mirrored?= =?UTF-8?q?=20from=20Q1=5F0=20not=20Q2=5F0=20=E2=80=94=20UNTESTED=20as=20c?= =?UTF-8?q?ommitted?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Adds the Vulkan path for PTQ1_0: the block struct and QUANT_K defines, the trit accessor plus dequantize/dequantize4/get_dm, the mul_mm load branch, a standalone dequant shader, the shader-generator registration, and 16 pipeline sites in ggml-vulkan.cpp mirrored from Q1_0's. Mirrored from Q1_0 deliberately. Q1_0 is type 41 at group 128 with correct Vulkan kernels; Q2_0 in this tree is group 128 in ggml-common.h but its Vulkan side still declares QUANT_K_Q2_0 64, because upstream's group-64 kernels were merged over a group-128 type and only some of the GLSL was re-derived. Pattern-matching off Q2_0 would have inherited that. Q1_0 also has no MMVQ entry, so this follows the same dequantize-based path rather than inventing one. One divergence from Q1_0 worth naming: PTQ1_0 puts the scale LAST in the block (qs[24], qh[2], then d) where Q1_0 puts d first. The GLSL struct mirrors the C field order exactly. Getting that wrong misindexes every block silently instead of failing, which is the same class of error as the Q2_0 breakage above. Not wired, and it is a real gap rather than an oversight: on-device quantization (cpy_f32_quant, set_rows, cpy_quant_f32). Those need a base-3 ENCODER in GLSL, which is a separate piece of work and is not needed to run a model -- weights are quantized offline by llama-quantize. **This has not been compiled or run.** There was no Vulkan toolchain on the machine it was written on. It is pushed so it can be built and tested on a GPU box, and it should be treated as unverified until test-backend-ops says otherwise. --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 16 ++++++ .../vulkan-shaders/dequant_funcs.glsl | 45 ++++++++++++++++ .../vulkan-shaders/dequant_ptq1_0.comp | 54 +++++++++++++++++++ .../vulkan-shaders/mul_mm_funcs.glsl | 15 ++++++ .../src/ggml-vulkan/vulkan-shaders/types.glsl | 21 ++++++++ .../vulkan-shaders/vulkan-shaders-gen.cpp | 3 +- 6 files changed, 153 insertions(+), 1 deletion(-) create mode 100644 ggml/src/ggml-vulkan/vulkan-shaders/dequant_ptq1_0.comp diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index 832e344ccffd..5a8e0a1a5257 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -4634,6 +4634,7 @@ 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) + CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_PTQ1_0], matmul_ptq1_0_f16, mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3) 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) @@ -4675,6 +4676,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { } #endif CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q1_0], matmul_id_subgroup_q1_0_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5) + CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_PTQ1_0], matmul_id_subgroup_ptq1_0_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5) CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_0], matmul_id_subgroup_q2_0_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5) CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_0], matmul_id_subgroup_q4_0_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5) CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_1], matmul_id_subgroup_q4_1_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5) @@ -4747,6 +4749,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 +4796,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 +4888,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 +4938,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 +4987,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 +5067,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 +5117,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 +5148,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 +5253,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 +5282,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 +5338,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 +5401,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 +5432,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 +5462,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); diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/dequant_funcs.glsl b/ggml/src/ggml-vulkan/vulkan-shaders/dequant_funcs.glsl index 627932bd3547..c9f847a57e58 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/dequant_funcs.glsl +++ b/ggml/src/ggml-vulkan/vulkan-shaders/dequant_funcs.glsl @@ -126,6 +126,44 @@ vec4 dequantize4(uint ib, uint iqs, uint a_offset) { } #endif +#if defined(DATA_A_PTQ1_0) +// Element order is not positional: a 16-byte chunk of qs, then an 8-byte chunk, then +// qh at four trits per byte, matching the CPU codec. 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); +} + +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 +602,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_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/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..c7f31a4aabd4 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"; From 1e29f207705fea3e77ac6d9944c816d55d6e75f0 Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Thu, 3 Sep 2026 00:51:01 -0700 Subject: [PATCH 11/14] =?UTF-8?q?cuda:=20PTQ1=5F0=20MMVQ=20takes=20the=20w?= =?UTF-8?q?hole=20block=20=E2=80=94=20measured,=20and=20still=206-9x=20slo?= =?UTF-8?q?wer?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Benchmarked the CUDA path for the first time. It is correct and slow, and the numbers matter more than the change: us/run on a 4090 at m=4096, k=14336, against pq2_0 on the same run. n=1 n=2 n=3 n=8 n=512 q1_0 8.88 9.82 13.05 30.14 225.4 pq2_0 9.75 11.05 14.92 31.27 232.4 ptq1_0 46.8 64.6 99.1 287.3 525.6 The per-chunk version asked "which byte holds element e" per element, a branch plus a variable-length recurrence, and iqs differs per lane so the branch diverged. Taking the whole 128-weight block per call (VDR 1 -> 4) makes every lane walk the same 26 bytes in the same order. That bought about 9% at n=2, 70.6 -> 64.6, and no more. So divergence was not the dominant cost, and the comment now says so. What remains is load shape: this issues 128 scalar int8 reads of the q8_1 activations, one per element, where pq2_0 issues eight dp4a calls over 4-byte words. The fix is identified but deliberately not guessed at. Below element 80 the element stride is 16 and 4 divides 16, so four consecutive elements are four consecutive BYTES at the same trit index. Four recurrences packed into one int would feed dp4a against a 4-byte q8_1 load, matching pq2_0's shape. That needs a build cycle to verify, and this session has already shown twice -- on the CPU vec_dot and on the Metal mapping -- that predictions about where GPU time goes are wrong about as often as they are right. Correctness held throughout: test-backend-ops -o MUL_MAT 1283/1283 on both an RTX 5090 and an RTX 4090, Backend CUDA0 OK, 45 ptq1_0 cases OK and 0 bad, before and after. Note n=512 is the cuBLAS dequantize path rather than MMVQ, and is 2.3x rather than 6-9x, which is consistent with the problem being in the MMVQ load pattern specifically. --- ggml/src/ggml-cuda/vecdotq.cuh | 83 +++++++++++++++++++++++++--------- 1 file changed, 62 insertions(+), 21 deletions(-) diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index 6484b15506ba..2915e0ab6725 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -113,7 +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 1 // one 32-element chunk at a time (4 chunks per 128 block) +#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 @@ -803,36 +803,77 @@ static __device__ __forceinline__ float vec_dot_q2_0_q8_1( return d2 * d8 * sumi; } -// PTQ1_0 x Q8_1. iqs selects one of the four 32-element chunks of the 128-weight -// block, pairing with one Q8_1 block, same convention as PQ2_0. Trits are gathered -// individually because base-3 packing gives no byte-aligned run inside a chunk, then -// packed four at a time into an int so the accumulation can still use dp4a. -// NOT COMPILED OR RUN: no CUDA toolchain was available where this was written. +// PTQ1_0 x Q8_1. One call consumes the whole 128-weight block and all four q8_1 +// blocks that cover it, so VDR is 4 rather than 1: that keeps the byte walk uniform +// across lanes instead of branching on which byte holds element e, which diverges +// because iqs differs per lane. +// +// MEASURED AND STILL SLOW. On a 4090 against pq2_0, us/run at m=4096 k=14336: +// n=1 46.8 -> (per-chunk version) pq2_0 9.75 +// n=2 70.6 -> 64.6 with this version pq2_0 11.05 +// n=8 287.3 pq2_0 31.27 +// So roughly 6x to 9x slower. Removing the divergent branch bought only ~9%, so the +// branch was not the dominant cost. The remaining cost is load shape: this issues 128 +// scalar int8 reads of the q8_1 activations, one per element, where pq2_0 issues eight +// dp4a calls over 4-byte words. +// +// The fix, not implemented here: for elements 4j..4j+3 below element 80 the stride is +// 16 and 4 divides 16, so four CONSECUTIVE elements are four consecutive BYTES at the +// same trit index. Four recurrences packed into one int would then feed dp4a against a +// 4-byte q8_1 load, matching pq2_0's load shape. That needs a build to verify and is +// left as the next step rather than guessed at. 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) { - const block_ptq1_0 * bq = (const block_ptq1_0 *) vbq + kbx; - const block_q8_1 * bq8 = bq8_1 + iqs; + const block_ptq1_0 * bq = (const block_ptq1_0 *) vbq + kbx; - const int base = iqs * 32; + int sumi[4] = {0, 0, 0, 0}; - int sumi = 0; #pragma unroll - 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); + for (int m = 0; m < 16; ++m) { // qs[0..15]: element m + 16*t + 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]; + } + } - // four signed int8 lanes, -1/0/1 - const int qx = (t0 & 0xFF) | ((t1 & 0xFF) << 8) | ((t2 & 0xFF) << 16) | ((t3 & 0xFF) << 24); - const int u = get_int_b4(bq8->qs, j); +#pragma unroll + for (int m = 0; m < 8; ++m) { // qs[16..23]: element 80 + 8*t + 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]; + } + } - sumi = ggml_cuda_dp4a(u, qx, sumi); +#pragma unroll + for (int h = 0; h < 2; ++h) { // qh[0..1]: element 120 + 2*t + 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]; + } } - const float d8 = __low2float(bq8->ds); - return (float) bq->d * d8 * sumi; + 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; } static __device__ __forceinline__ float vec_dot_pq2_0_q8_1( From 599d0b8768edf99c91586977495a46a92198a3d0 Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Thu, 3 Sep 2026 07:47:41 -0700 Subject: [PATCH 12/14] =?UTF-8?q?vulkan:=20fix=20three=20build=20breaks=20?= =?UTF-8?q?=E2=80=94=20shared=20ptq1=5F0.glsl,=20skip=20coopmat2,=20drop?= =?UTF-8?q?=20cm2=20pipelines?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The Vulkan commit did not build. Three errors, one root cause each, all found by a teammate on an Intel Arc B390 and independently by windows CI: 1. mul_mm_cm2.comp: 'dequantFuncA' undeclared. Adding the type to type_names in vulkan-shaders-gen.cpp auto-enrols it in coopmat2 generation, but dequant_funcs_cm2.glsl has no PTQ1_0 decoder (Q1_0 has six entries). That failed the whole Vulkan shader build, not just this type. Now skipped for ptq1_0, which falls back to the scalar and coopmat1 paths that are actually implemented. 2. mul_mm_funcs.glsl: 'ptq1_0_trit' no matching function. The helper lived in dequant_funcs.glsl; mul_mm.comp includes types/dot_product/mul_mm_id/mul_mm_funcs and never dequant_funcs.glsl. **This is the same bug as the CUDA one fixed an hour earlier, same helper, different backend**: defined in the dequantize chain, called from the matmul chain. Fixed the same way, with a shared header vulkan-shaders/ptq1_0.glsl behind an include guard, included from both. Not duplicated: two copies would have to stay in step with the CPU codec in ggml-quants.c, and a divergence there appears as wrong matmul results, never as a build error. 3. ggml-vulkan.cpp: 'matmul_ptq1_0_f16_cm2_len' undeclared. Registrations for cm2 pipelines whose shaders no longer exist once (1) is fixed. Both removed, with a comment recording that the absence is deliberate. (1) and (3) had to be fixed together or one error just trades for the other. Verified locally this time, which is the other half of the lesson: glslc, molten-vk, vulkan-headers and vulkan-loader were already installed on the machine that wrote the original commit. It was pushed unverified on the false assumption that no toolchain was present, and a local build finds all three in minutes. -DGGML_VULKAN=ON builds clean SPIR-V generated 1991 shaders, 54 of them ptq1_0, incl. coopmat1, no cm2 test-backend-ops MUL_MAT 1009/1009, Backend Vulkan0: OK, no regression That local run does NOT verify the kernel. MoltenVK reports every ptq1_0 case "not supported", and it reports pq2_0 the same way (q1_0 gets 28 of 156), so this is an Apple/MoltenVK coverage limit rather than a defect in the new type. Kernel correctness comes from the Arc B390 run on the patched tree: 1009/1009 with 78 ptq1_0 cases actually executed and 0 failures, on KHR_coopmat. Still absent by design: MMVQ (no q8_1 pipelines, absent from should_use_mmvq), coopmat2, and on-device quantization, which needs a base-3 encoder in GLSL. --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 5 ++- .../vulkan-shaders/dequant_funcs.glsl | 26 +------------ .../ggml-vulkan/vulkan-shaders/mul_mm.comp | 3 ++ .../ggml-vulkan/vulkan-shaders/ptq1_0.glsl | 38 +++++++++++++++++++ .../vulkan-shaders/vulkan-shaders-gen.cpp | 7 ++++ 5 files changed, 52 insertions(+), 27 deletions(-) create mode 100644 ggml/src/ggml-vulkan/vulkan-shaders/ptq1_0.glsl diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index 5a8e0a1a5257..e80056009b4a 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -4634,7 +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) - CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_PTQ1_0], matmul_ptq1_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) @@ -4676,7 +4678,6 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) { } #endif CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q1_0], matmul_id_subgroup_q1_0_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5) - CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_PTQ1_0], matmul_id_subgroup_ptq1_0_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5) CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q2_0], matmul_id_subgroup_q2_0_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5) CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_0], matmul_id_subgroup_q4_0_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5) CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_Q4_1], matmul_id_subgroup_q4_1_f16, mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5) diff --git a/ggml/src/ggml-vulkan/vulkan-shaders/dequant_funcs.glsl b/ggml/src/ggml-vulkan/vulkan-shaders/dequant_funcs.glsl index c9f847a57e58..88be1360403a 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/dequant_funcs.glsl +++ b/ggml/src/ggml-vulkan/vulkan-shaders/dequant_funcs.glsl @@ -127,31 +127,7 @@ vec4 dequantize4(uint ib, uint iqs, uint a_offset) { #endif #if defined(DATA_A_PTQ1_0) -// Element order is not positional: a 16-byte chunk of qs, then an 8-byte chunk, then -// qh at four trits per byte, matching the CPU codec. 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); -} +#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)); 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/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/vulkan-shaders-gen.cpp b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp index c7f31a4aabd4..83af8af068fe 100644 --- a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp +++ b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp @@ -592,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 From 705d9ccda8fb4dc8367a35eea6331c238204e877 Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Thu, 3 Sep 2026 08:52:28 -0700 Subject: [PATCH 13/14] vulkan: make PTQ1_0 reachable, and guard the coopmat2 path The PTQ1_0 pipelines were created but the type was absent from every selection and support enumeration, so supports_op returned false, the scheduler never offered MUL_MAT/MUL_MAT_ID/GET_ROWS to Vulkan, and the generated shaders were unreachable. The build looked green while nothing executed on the backend. Add the type alongside GGML_TYPE_Q1_0 in the five pipeline getters and the two supports_op gates. Adding it there alone is not safe, because PTQ1_0 is the only type in those lists with no coopmat2 pipelines: dequant_funcs_cm2.glsl carries no decoder for it, so vulkan-shaders-gen skips cm2 generation. The cm2 slots are still allocated by vk_matmul_pipeline2's constructor, so the mm getter would return a non-null pipeline whose l/m/s are all null, and the caller's "mmp == nullptr" fallback would not fire; the mul_mat_id getter would instead trip GGML_ASSERT(support_fp32acc). Both getters now return nullptr when the cm2 slots are empty, which routes the caller to the dequant + f16 matmul fallback it already implements. That path needs ggml_vk_get_to_fp16, which is one of the gates opened above. The cpy, SET_ROWS and DUP gates are deliberately left alone: copy_to_quant.comp has no PTQ1_0 block, so quantizing to this type on device is unimplemented and a refusal there is correct. Verified on Intel Arc B390 (KHR_coopmat): ptq1_0 goes from 0 OK / 78 declined to 28 OK / 50 declined / 0 failures, matching q1_0 and q2_0 case for case, byte-identical across two runs. The coopmat2 guard is reasoned from the pipeline-creation structure and is not covered by that run, which is a coopmat1 device; it has no coopmat2 hardware behind it yet. --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 24 +++++++++++++++++++++++- 1 file changed, 23 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index e80056009b4a..e30fb0d63329 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -7673,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: @@ -7748,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: @@ -7778,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; @@ -7818,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: @@ -7912,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: @@ -7946,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 { @@ -7985,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: @@ -18165,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: @@ -18271,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: From d213a381ba32ef5d4663a0f536e5bb3bb7b5a09d Mon Sep 17 00:00:00 2001 From: bri-prism <288398250+bri-prism@users.noreply.github.com> Date: Thu, 3 Sep 2026 12:40:01 -0700 Subject: [PATCH 14/14] cuda: accelerate PTQ1_0 MMVQ with packed DP4A Assisted-by: OpenAI Codex --- ggml/src/ggml-cuda/getrows.cu | 1 + ggml/src/ggml-cuda/mmvq.cu | 46 +++++++-- ggml/src/ggml-cuda/vecdotq.cuh | 166 ++++++++++++++++++++++++--------- gguf-py/gguf/constants.py | 9 +- 4 files changed, 165 insertions(+), 57 deletions(-) diff --git a/ggml/src/ggml-cuda/getrows.cu b/ggml/src/ggml-cuda/getrows.cu index c3f365458df2..6b4453bc69ce 100644 --- a/ggml/src/ggml-cuda/getrows.cu +++ b/ggml/src/ggml-cuda/getrows.cu @@ -323,6 +323,7 @@ static void ggml_cuda_get_rows_switch_src0_type( case GGML_TYPE_PQ2_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_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); diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 5ac3eaf30c4d..ede8a65e0154 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -665,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 diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index 2915e0ab6725..fa6cdabd5d84 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -803,82 +803,158 @@ static __device__ __forceinline__ float vec_dot_q2_0_q8_1( return d2 * d8 * sumi; } -// PTQ1_0 x Q8_1. One call consumes the whole 128-weight block and all four q8_1 -// blocks that cover it, so VDR is 4 rather than 1: that keeps the byte walk uniform -// across lanes instead of branching on which byte holds element e, which diverges -// because iqs differs per lane. -// -// MEASURED AND STILL SLOW. On a 4090 against pq2_0, us/run at m=4096 k=14336: -// n=1 46.8 -> (per-chunk version) pq2_0 9.75 -// n=2 70.6 -> 64.6 with this version pq2_0 11.05 -// n=8 287.3 pq2_0 31.27 -// So roughly 6x to 9x slower. Removing the divergent branch bought only ~9%, so the -// branch was not the dominant cost. The remaining cost is load shape: this issues 128 -// scalar int8 reads of the q8_1 activations, one per element, where pq2_0 issues eight -// dp4a calls over 4-byte words. -// -// The fix, not implemented here: for elements 4j..4j+3 below element 80 the stride is -// 16 and 4 divides 16, so four CONSECUTIVE elements are four consecutive BYTES at the -// same trit index. Four recurrences packed into one int would then feed dp4a against a -// 4-byte q8_1 load, matching pq2_0's load shape. That needs a build to verify and is -// left as the next step rather than guessed at. -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) +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]); + } + } + } - const block_ptq1_0 * bq = (const block_ptq1_0 *) vbq + kbx; +# 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); - int sumi[4] = {0, 0, 0, 0}; +# 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]); + } + } + } -#pragma unroll - for (int m = 0; m < 16; ++m) { // qs[0..15]: element m + 16*t + 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 +# 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; + 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) { // qs[16..23]: element 80 + 8*t + m +# pragma unroll + for (int m = 0; m < 8; ++m) { uint32_t v = bq->qs[16 + m]; -#pragma unroll +# 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; + 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) { // qh[0..1]: element 120 + 2*t + h +# pragma unroll + for (int h = 0; h < 2; ++h) { uint32_t v = bq->qh[h]; -#pragma unroll +# 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; + 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 +# 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) { - +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/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), }