diff --git a/include/ggml.h b/include/ggml.h index d6807b6d..01a8899e 100644 --- a/include/ggml.h +++ b/include/ggml.h @@ -429,7 +429,27 @@ extern "C" { GGML_TYPE_MXFP4 = 39, // MXFP4 (1 block) GGML_TYPE_NVFP4 = 40, // NVFP4 (4 blocks, E4M3 scale) GGML_TYPE_Q1_0 = 41, - GGML_TYPE_COUNT = 42, + GGML_TYPE_TQ2P = 42, // sasori dual-trit-plane (two trit planes + two per-block scales) + GGML_TYPE_TQ3P = 43, // sasori tri-trit-plane (three trit planes + three per-block scales) + GGML_TYPE_TQ1P = 44, // sasori single-trit-plane (one trit plane + one per-block scale) + GGML_TYPE_TQ4P = 45, // sasori quad-trit-plane (four trit planes + four per-block scales) + GGML_TYPE_TQ1P_G32 = 46, // sasori K=1 g=32 + GGML_TYPE_TQ1P_G64 = 47, // sasori K=1 g=64 + GGML_TYPE_TQ1P_G128 = 48, // sasori K=1 g=128 + GGML_TYPE_TQ2P_G32 = 49, // sasori K=2 g=32 + GGML_TYPE_TQ2P_G64 = 50, // sasori K=2 g=64 + GGML_TYPE_TQ2P_G128 = 51, // sasori K=2 g=128 + GGML_TYPE_TQ3P_G32 = 52, // sasori K=3 g=32 + GGML_TYPE_TQ3P_G64 = 53, // sasori K=3 g=64 + GGML_TYPE_TQ3P_G128 = 54, // sasori K=3 g=128 + GGML_TYPE_TQ4P_G32 = 55, // sasori K=4 g=32 + GGML_TYPE_TQ4P_G64 = 56, // sasori K=4 g=64 + GGML_TYPE_TQ4P_G128 = 57, // sasori K=4 g=128 + GGML_TYPE_TQ1P_Q4K6 = 58, // sasori K=1 g=32, q4k6 scale (6-bit sub + fp16 super/256) + GGML_TYPE_TQ2P_Q4K6 = 59, // sasori K=2 g=32, q4k6 scale (BYTE_PP, 4.625 bpw) + GGML_TYPE_TQ3P_Q4K6 = 60, // sasori K=3 g=32, q4k6 scale + GGML_TYPE_TQ1PA = 61, // sasori K=1 asym (trit plane + f16 scale + f16 mu offset) + GGML_TYPE_COUNT = 62, }; // precision diff --git a/src/ggml-common.h b/src/ggml-common.h index f05683b4..13531c8b 100644 --- a/src/ggml-common.h +++ b/src/ggml-common.h @@ -277,6 +277,68 @@ typedef struct { } block_tq2_0; static_assert(sizeof(block_tq2_0) == sizeof(ggml_half) + QK_K / 4, "wrong tq2_0 block size/padding"); +// sasori dual-trit-plane: two trit planes (2 bits/trit) + two per-block fp16 scales. ~4.125 bpw. +// Block size and the strided 2-bit layout mirror block_tq2_0 (per plane) so the AVX2 vec_dot +// can pair each extraction with a contiguous q8_K activation chunk (see ggml_vec_dot_tq2p_q8_K). +#define QK_TQ2P 256 +typedef struct { + uint8_t qs1[QK_TQ2P/4]; // plane 1, 2 bits per trit (tq2_0 strided layout); trit = (2-bit value) - 1 + uint8_t qs2[QK_TQ2P/4]; // plane 2 + ggml_half d1; // plane-1 scale + ggml_half d2; // plane-2 scale +} block_tq2p; +static_assert(sizeof(block_tq2p) == 2*(QK_TQ2P/4) + 2*sizeof(ggml_half), "wrong tq2p block size/padding"); + +#define QK_TQ3P 256 +typedef struct { + uint8_t qs1[QK_TQ3P/4]; // plane 1, 2 bits per trit (tq2_0 strided layout); trit = (2-bit value) - 1 + uint8_t qs2[QK_TQ3P/4]; // plane 2 + uint8_t qs3[QK_TQ3P/4]; // plane 3 (the uniform-K=3 type; mirrors block_tq2p + one extra plane) + ggml_half d1; // plane-1 scale + ggml_half d2; // plane-2 scale + ggml_half d3; // plane-3 scale +} block_tq3p; +static_assert(sizeof(block_tq3p) == 3*(QK_TQ3P/4) + 3*sizeof(ggml_half), "wrong tq3p block size/padding"); + +#define QK_TQ1P 256 +typedef struct { + uint8_t qs1[QK_TQ1P/4]; // single plane, 2 bits per trit (tq2_0 strided layout) + ggml_half d; // plane scale +} block_tq1p; +static_assert(sizeof(block_tq1p) == (QK_TQ1P/4) + sizeof(ggml_half), "wrong tq1p block size/padding"); + +// sasori K=1 assimétrico (asym-μ): = block_tq1p + um offset μ f16 por grupo (g=256, nsub=1). +// Decode por peso: wh = d*trit + μ. μ vem DEPOIS da escala (layout | qs1 | d | μ |). 68 B = 2.125 bpw. +typedef struct { + uint8_t qs1[QK_TQ1P/4]; // plano único, 2 bits por trit (layout strided tq2_0) — igual ao block_tq1p + ggml_half d; // escala do grupo (g256) + ggml_half mu; // offset μ do grupo (g256), somado a todo peso do super-bloco +} block_tq1pa; +static_assert(sizeof(block_tq1pa) == (QK_TQ1P/4) + 2*sizeof(ggml_half), "wrong tq1pa block size/padding"); + +#define QK_TQ4P 256 +typedef struct { + uint8_t qs1[QK_TQ4P/4]; // plane 1, 2 bits per trit (tq2_0 strided layout) + uint8_t qs2[QK_TQ4P/4]; // plane 2 + uint8_t qs3[QK_TQ4P/4]; // plane 3 + uint8_t qs4[QK_TQ4P/4]; // plane 4 (= block_tq3p + one extra plane) + ggml_half d1; + ggml_half d2; + ggml_half d3; + ggml_half d4; +} block_tq4p; +static_assert(sizeof(block_tq4p) == 4*(QK_TQ4P/4) + 4*sizeof(ggml_half), "wrong tq4p block size/padding"); + +// sasori g-variable types: 256-trit strided super-block per plane + nsub = 256/g f16 scales/plane; +// no named struct, type_size = TQKP_GV_BYTES(K, g). +#define QK_TQKP_SUPER 256 +#define TQKP_GV_BYTES(K, g) ((K)*(QK_TQKP_SUPER/4) + (K)*(QK_TQKP_SUPER/(g))*2) + +// sasori q4k6 scale format (g=32 only): K 256-strided planes (same as TQKP_GV) + K*nsub 1-byte +// sub-scales (bit7 sign, bits5..0 u6) + K fp16 super-scales/256 (BYTE_PP). nsub = 256/32 = 8. +#define Q4K6_G 32 +#define TQKP_Q4K6_BYTES(K) ((K)*(QK_TQKP_SUPER/4) + (K)*(QK_TQKP_SUPER/Q4K6_G) + (K)*2) + // // Super-block quantization structures // diff --git a/src/ggml-cpu/arch/x86/quants.c b/src/ggml-cpu/arch/x86/quants.c index 94b19b82..7b219ae7 100644 --- a/src/ggml-cpu/arch/x86/quants.c +++ b/src/ggml-cpu/arch/x86/quants.c @@ -1433,6 +1433,576 @@ void ggml_vec_dot_tq2_0_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const vo #endif } +void ggml_vec_dot_tq2p_q8_K(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) { + assert(n % QK_TQ2P == 0); + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + + const block_tq2p * GGML_RESTRICT x = vx; + const block_q8_K * GGML_RESTRICT y = vy; + + const int nb = n / QK_TQ2P; // QK_TQ2P == QK_K -> 1:1 with q8_K blocks + +#if defined(__AVX2__) + // Two trit planes share the SAME q8_K activations: load qy once, maddubs twice (TQ2_0 trick). + // u = trit + 1 in {0,1,2}; sum(u*y) - sum(y) = sum(trit*y). The ysum (=q8_K bsums) subtraction + // happens before the horizontal reduction, so per-lane scrambling is irrelevant to the total. + __m256 sumf = _mm256_setzero_ps(); + + for (int i = 0; i < nb; ++i) { + __m256i a0 = _mm256_setzero_si256(); // plane 1 + __m256i a1 = _mm256_setzero_si256(); + __m256i b0 = _mm256_setzero_si256(); // plane 2 + __m256i b1 = _mm256_setzero_si256(); + + for (size_t j = 0; j < sizeof(x->qs1); j += 32) { + const __m256i qy0 = _mm256_loadu_si256((const __m256i *) (y[i].qs + j*4 + 0)); + const __m256i qy1 = _mm256_loadu_si256((const __m256i *) (y[i].qs + j*4 + 32)); + const __m256i qy2 = _mm256_loadu_si256((const __m256i *) (y[i].qs + j*4 + 64)); + const __m256i qy3 = _mm256_loadu_si256((const __m256i *) (y[i].qs + j*4 + 96)); + + __m256i p0 = _mm256_loadu_si256((const __m256i *) (x[i].qs1 + j)); + __m256i p1 = _mm256_and_si256(_mm256_srli_epi16(p0, 2), _mm256_set1_epi8(3)); + __m256i p2 = _mm256_and_si256(_mm256_srli_epi16(p0, 4), _mm256_set1_epi8(3)); + __m256i p3 = _mm256_and_si256(_mm256_srli_epi16(p0, 6), _mm256_set1_epi8(3)); + p0 = _mm256_and_si256(p0, _mm256_set1_epi8(3)); + a0 = _mm256_add_epi16(a0, _mm256_add_epi16(_mm256_maddubs_epi16(p0, qy0), _mm256_maddubs_epi16(p1, qy1))); + a1 = _mm256_add_epi16(a1, _mm256_add_epi16(_mm256_maddubs_epi16(p2, qy2), _mm256_maddubs_epi16(p3, qy3))); + + __m256i r0 = _mm256_loadu_si256((const __m256i *) (x[i].qs2 + j)); + __m256i r1 = _mm256_and_si256(_mm256_srli_epi16(r0, 2), _mm256_set1_epi8(3)); + __m256i r2 = _mm256_and_si256(_mm256_srli_epi16(r0, 4), _mm256_set1_epi8(3)); + __m256i r3 = _mm256_and_si256(_mm256_srli_epi16(r0, 6), _mm256_set1_epi8(3)); + r0 = _mm256_and_si256(r0, _mm256_set1_epi8(3)); + b0 = _mm256_add_epi16(b0, _mm256_add_epi16(_mm256_maddubs_epi16(r0, qy0), _mm256_maddubs_epi16(r1, qy1))); + b1 = _mm256_add_epi16(b1, _mm256_add_epi16(_mm256_maddubs_epi16(r2, qy2), _mm256_maddubs_epi16(r3, qy3))); + } + + const __m256i ysum = _mm256_loadu_si256((const __m256i *) y[i].bsums); + const __m256 d1 = _mm256_set1_ps(y[i].d * GGML_CPU_FP16_TO_FP32(x[i].d1)); + const __m256 d2 = _mm256_set1_ps(y[i].d * GGML_CPU_FP16_TO_FP32(x[i].d2)); + + __m256i sa = _mm256_madd_epi16(_mm256_sub_epi16(_mm256_add_epi16(a0, a1), ysum), _mm256_set1_epi16(1)); + __m256i sb = _mm256_madd_epi16(_mm256_sub_epi16(_mm256_add_epi16(b0, b1), ysum), _mm256_set1_epi16(1)); + + sumf = _mm256_add_ps(_mm256_mul_ps(_mm256_cvtepi32_ps(sa), d1), sumf); + sumf = _mm256_add_ps(_mm256_mul_ps(_mm256_cvtepi32_ps(sb), d2), sumf); + } + + *s = hsum_float_8(sumf); +#else + UNUSED(x); + UNUSED(y); + UNUSED(nb); + ggml_vec_dot_tq2p_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} + +void ggml_vec_dot_tq3p_q8_K(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) { + assert(n % QK_TQ3P == 0); + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + + const block_tq3p * GGML_RESTRICT x = vx; + const block_q8_K * GGML_RESTRICT y = vy; + + const int nb = n / QK_TQ3P; // QK_TQ3P == QK_K -> 1:1 with q8_K blocks + +#if defined(__AVX2__) + // Three trit planes share the SAME q8_K activations: load qy once, maddubs three times (one per + // plane). Same TQ2_0 trick/ysum subtraction as TQ2P; per-block int16 accumulators stay well + // within range (|(a0+a1)-ysum| <~ 6096 << 32767) and cross-block accumulation is in fp32. + __m256 sumf = _mm256_setzero_ps(); + + for (int i = 0; i < nb; ++i) { + __m256i a0 = _mm256_setzero_si256(); // plane 1 + __m256i a1 = _mm256_setzero_si256(); + __m256i b0 = _mm256_setzero_si256(); // plane 2 + __m256i b1 = _mm256_setzero_si256(); + __m256i c0 = _mm256_setzero_si256(); // plane 3 + __m256i c1 = _mm256_setzero_si256(); + + for (size_t j = 0; j < sizeof(x->qs1); j += 32) { + const __m256i qy0 = _mm256_loadu_si256((const __m256i *) (y[i].qs + j*4 + 0)); + const __m256i qy1 = _mm256_loadu_si256((const __m256i *) (y[i].qs + j*4 + 32)); + const __m256i qy2 = _mm256_loadu_si256((const __m256i *) (y[i].qs + j*4 + 64)); + const __m256i qy3 = _mm256_loadu_si256((const __m256i *) (y[i].qs + j*4 + 96)); + + __m256i p0 = _mm256_loadu_si256((const __m256i *) (x[i].qs1 + j)); + __m256i p1 = _mm256_and_si256(_mm256_srli_epi16(p0, 2), _mm256_set1_epi8(3)); + __m256i p2 = _mm256_and_si256(_mm256_srli_epi16(p0, 4), _mm256_set1_epi8(3)); + __m256i p3 = _mm256_and_si256(_mm256_srli_epi16(p0, 6), _mm256_set1_epi8(3)); + p0 = _mm256_and_si256(p0, _mm256_set1_epi8(3)); + a0 = _mm256_add_epi16(a0, _mm256_add_epi16(_mm256_maddubs_epi16(p0, qy0), _mm256_maddubs_epi16(p1, qy1))); + a1 = _mm256_add_epi16(a1, _mm256_add_epi16(_mm256_maddubs_epi16(p2, qy2), _mm256_maddubs_epi16(p3, qy3))); + + __m256i r0 = _mm256_loadu_si256((const __m256i *) (x[i].qs2 + j)); + __m256i r1 = _mm256_and_si256(_mm256_srli_epi16(r0, 2), _mm256_set1_epi8(3)); + __m256i r2 = _mm256_and_si256(_mm256_srli_epi16(r0, 4), _mm256_set1_epi8(3)); + __m256i r3 = _mm256_and_si256(_mm256_srli_epi16(r0, 6), _mm256_set1_epi8(3)); + r0 = _mm256_and_si256(r0, _mm256_set1_epi8(3)); + b0 = _mm256_add_epi16(b0, _mm256_add_epi16(_mm256_maddubs_epi16(r0, qy0), _mm256_maddubs_epi16(r1, qy1))); + b1 = _mm256_add_epi16(b1, _mm256_add_epi16(_mm256_maddubs_epi16(r2, qy2), _mm256_maddubs_epi16(r3, qy3))); + + __m256i w0 = _mm256_loadu_si256((const __m256i *) (x[i].qs3 + j)); + __m256i w1 = _mm256_and_si256(_mm256_srli_epi16(w0, 2), _mm256_set1_epi8(3)); + __m256i w2 = _mm256_and_si256(_mm256_srli_epi16(w0, 4), _mm256_set1_epi8(3)); + __m256i w3 = _mm256_and_si256(_mm256_srli_epi16(w0, 6), _mm256_set1_epi8(3)); + w0 = _mm256_and_si256(w0, _mm256_set1_epi8(3)); + c0 = _mm256_add_epi16(c0, _mm256_add_epi16(_mm256_maddubs_epi16(w0, qy0), _mm256_maddubs_epi16(w1, qy1))); + c1 = _mm256_add_epi16(c1, _mm256_add_epi16(_mm256_maddubs_epi16(w2, qy2), _mm256_maddubs_epi16(w3, qy3))); + } + + const __m256i ysum = _mm256_loadu_si256((const __m256i *) y[i].bsums); + const __m256 d1 = _mm256_set1_ps(y[i].d * GGML_CPU_FP16_TO_FP32(x[i].d1)); + const __m256 d2 = _mm256_set1_ps(y[i].d * GGML_CPU_FP16_TO_FP32(x[i].d2)); + const __m256 d3 = _mm256_set1_ps(y[i].d * GGML_CPU_FP16_TO_FP32(x[i].d3)); + + __m256i sa = _mm256_madd_epi16(_mm256_sub_epi16(_mm256_add_epi16(a0, a1), ysum), _mm256_set1_epi16(1)); + __m256i sb = _mm256_madd_epi16(_mm256_sub_epi16(_mm256_add_epi16(b0, b1), ysum), _mm256_set1_epi16(1)); + __m256i sc = _mm256_madd_epi16(_mm256_sub_epi16(_mm256_add_epi16(c0, c1), ysum), _mm256_set1_epi16(1)); + + sumf = _mm256_add_ps(_mm256_mul_ps(_mm256_cvtepi32_ps(sa), d1), sumf); + sumf = _mm256_add_ps(_mm256_mul_ps(_mm256_cvtepi32_ps(sb), d2), sumf); + sumf = _mm256_add_ps(_mm256_mul_ps(_mm256_cvtepi32_ps(sc), d3), sumf); + } + + *s = hsum_float_8(sumf); +#else + UNUSED(x); + UNUSED(y); + UNUSED(nb); + ggml_vec_dot_tq3p_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} + +void ggml_vec_dot_tq1p_q8_K(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) { + assert(n % QK_TQ1P == 0); + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + + const block_tq1p * GGML_RESTRICT x = vx; + const block_q8_K * GGML_RESTRICT y = vy; + + const int nb = n / QK_TQ1P; + +#if defined(__AVX2__) + // One trit plane; same TQ2_0 maddubs/ysum scheme as tq2p (one plane instead of two). + __m256 sumf = _mm256_setzero_ps(); + + for (int i = 0; i < nb; ++i) { + __m256i a0 = _mm256_setzero_si256(); + __m256i a1 = _mm256_setzero_si256(); + + for (size_t j = 0; j < sizeof(x->qs1); j += 32) { + const __m256i qy0 = _mm256_loadu_si256((const __m256i *) (y[i].qs + j*4 + 0)); + const __m256i qy1 = _mm256_loadu_si256((const __m256i *) (y[i].qs + j*4 + 32)); + const __m256i qy2 = _mm256_loadu_si256((const __m256i *) (y[i].qs + j*4 + 64)); + const __m256i qy3 = _mm256_loadu_si256((const __m256i *) (y[i].qs + j*4 + 96)); + + __m256i p0 = _mm256_loadu_si256((const __m256i *) (x[i].qs1 + j)); + __m256i p1 = _mm256_and_si256(_mm256_srli_epi16(p0, 2), _mm256_set1_epi8(3)); + __m256i p2 = _mm256_and_si256(_mm256_srli_epi16(p0, 4), _mm256_set1_epi8(3)); + __m256i p3 = _mm256_and_si256(_mm256_srli_epi16(p0, 6), _mm256_set1_epi8(3)); + p0 = _mm256_and_si256(p0, _mm256_set1_epi8(3)); + a0 = _mm256_add_epi16(a0, _mm256_add_epi16(_mm256_maddubs_epi16(p0, qy0), _mm256_maddubs_epi16(p1, qy1))); + a1 = _mm256_add_epi16(a1, _mm256_add_epi16(_mm256_maddubs_epi16(p2, qy2), _mm256_maddubs_epi16(p3, qy3))); + } + + const __m256i ysum = _mm256_loadu_si256((const __m256i *) y[i].bsums); + const __m256 d1 = _mm256_set1_ps(y[i].d * GGML_CPU_FP16_TO_FP32(x[i].d)); + + __m256i sa = _mm256_madd_epi16(_mm256_sub_epi16(_mm256_add_epi16(a0, a1), ysum), _mm256_set1_epi16(1)); + + sumf = _mm256_add_ps(_mm256_mul_ps(_mm256_cvtepi32_ps(sa), d1), sumf); + } + + *s = hsum_float_8(sumf); +#else + UNUSED(x); + UNUSED(y); + UNUSED(nb); + ggml_vec_dot_tq1p_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} + +void ggml_vec_dot_tq1pa_q8_K(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) { + assert(n % QK_TQ1P == 0); + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + + const block_tq1pa * GGML_RESTRICT x = vx; + const block_q8_K * GGML_RESTRICT y = vy; + + const int nb = n / QK_TQ1P; + +#if defined(__AVX2__) + // sasori K=1 assimétrico: mesmo esquema maddubs/ysum do tq1p (um plano) + o termo mu*ysum do grupo. + // wh = d*trit + mu e (a0+a1)-ysum = Σ trit*yq => dot = d8*( d*(su1-ysum) + mu*ysum ). + // g256/nsub=1: o ysum do grupo == o ysum do bloco (Σ bsums), então mu entra SEM laço novo. + __m256 sumf = _mm256_setzero_ps(); + + for (int i = 0; i < nb; ++i) { + __m256i a0 = _mm256_setzero_si256(); + __m256i a1 = _mm256_setzero_si256(); + + for (size_t j = 0; j < sizeof(x->qs1); j += 32) { + const __m256i qy0 = _mm256_loadu_si256((const __m256i *) (y[i].qs + j*4 + 0)); + const __m256i qy1 = _mm256_loadu_si256((const __m256i *) (y[i].qs + j*4 + 32)); + const __m256i qy2 = _mm256_loadu_si256((const __m256i *) (y[i].qs + j*4 + 64)); + const __m256i qy3 = _mm256_loadu_si256((const __m256i *) (y[i].qs + j*4 + 96)); + + __m256i p0 = _mm256_loadu_si256((const __m256i *) (x[i].qs1 + j)); + __m256i p1 = _mm256_and_si256(_mm256_srli_epi16(p0, 2), _mm256_set1_epi8(3)); + __m256i p2 = _mm256_and_si256(_mm256_srli_epi16(p0, 4), _mm256_set1_epi8(3)); + __m256i p3 = _mm256_and_si256(_mm256_srli_epi16(p0, 6), _mm256_set1_epi8(3)); + p0 = _mm256_and_si256(p0, _mm256_set1_epi8(3)); + a0 = _mm256_add_epi16(a0, _mm256_add_epi16(_mm256_maddubs_epi16(p0, qy0), _mm256_maddubs_epi16(p1, qy1))); + a1 = _mm256_add_epi16(a1, _mm256_add_epi16(_mm256_maddubs_epi16(p2, qy2), _mm256_maddubs_epi16(p3, qy3))); + } + + const __m256i ysum = _mm256_loadu_si256((const __m256i *) y[i].bsums); + const __m256 d1 = _mm256_set1_ps(y[i].d * GGML_CPU_FP16_TO_FP32(x[i].d)); + const __m256 dmu = _mm256_set1_ps(y[i].d * GGML_CPU_FP16_TO_FP32(x[i].mu)); + + __m256i sa = _mm256_madd_epi16(_mm256_sub_epi16(_mm256_add_epi16(a0, a1), ysum), _mm256_set1_epi16(1)); + // asym: madd(ysum,1) espalha Σ bsums (= Σ yq do bloco, g256) em 8 int32; entra no MESMO hsum. + __m256i sy = _mm256_madd_epi16(ysum, _mm256_set1_epi16(1)); + + sumf = _mm256_add_ps(_mm256_mul_ps(_mm256_cvtepi32_ps(sa), d1), sumf); + sumf = _mm256_add_ps(_mm256_mul_ps(_mm256_cvtepi32_ps(sy), dmu), sumf); + } + + *s = hsum_float_8(sumf); +#else + UNUSED(x); + UNUSED(y); + UNUSED(nb); + ggml_vec_dot_tq1pa_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} + +void ggml_vec_dot_tq4p_q8_K(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) { + assert(n % QK_TQ4P == 0); + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + + const block_tq4p * GGML_RESTRICT x = vx; + const block_q8_K * GGML_RESTRICT y = vy; + + const int nb = n / QK_TQ4P; + +#if defined(__AVX2__) + // Four trit planes; same scheme as tq3p + one plane. Per-plane int16 accumulators stay in range + // (independent per plane), cross-block accumulation is fp32. + __m256 sumf = _mm256_setzero_ps(); + + for (int i = 0; i < nb; ++i) { + __m256i a0 = _mm256_setzero_si256(); // plane 1 + __m256i a1 = _mm256_setzero_si256(); + __m256i b0 = _mm256_setzero_si256(); // plane 2 + __m256i b1 = _mm256_setzero_si256(); + __m256i c0 = _mm256_setzero_si256(); // plane 3 + __m256i c1 = _mm256_setzero_si256(); + __m256i e0 = _mm256_setzero_si256(); // plane 4 + __m256i e1 = _mm256_setzero_si256(); + + for (size_t j = 0; j < sizeof(x->qs1); j += 32) { + const __m256i qy0 = _mm256_loadu_si256((const __m256i *) (y[i].qs + j*4 + 0)); + const __m256i qy1 = _mm256_loadu_si256((const __m256i *) (y[i].qs + j*4 + 32)); + const __m256i qy2 = _mm256_loadu_si256((const __m256i *) (y[i].qs + j*4 + 64)); + const __m256i qy3 = _mm256_loadu_si256((const __m256i *) (y[i].qs + j*4 + 96)); + + __m256i p0 = _mm256_loadu_si256((const __m256i *) (x[i].qs1 + j)); + __m256i p1 = _mm256_and_si256(_mm256_srli_epi16(p0, 2), _mm256_set1_epi8(3)); + __m256i p2 = _mm256_and_si256(_mm256_srli_epi16(p0, 4), _mm256_set1_epi8(3)); + __m256i p3 = _mm256_and_si256(_mm256_srli_epi16(p0, 6), _mm256_set1_epi8(3)); + p0 = _mm256_and_si256(p0, _mm256_set1_epi8(3)); + a0 = _mm256_add_epi16(a0, _mm256_add_epi16(_mm256_maddubs_epi16(p0, qy0), _mm256_maddubs_epi16(p1, qy1))); + a1 = _mm256_add_epi16(a1, _mm256_add_epi16(_mm256_maddubs_epi16(p2, qy2), _mm256_maddubs_epi16(p3, qy3))); + + __m256i r0 = _mm256_loadu_si256((const __m256i *) (x[i].qs2 + j)); + __m256i r1 = _mm256_and_si256(_mm256_srli_epi16(r0, 2), _mm256_set1_epi8(3)); + __m256i r2 = _mm256_and_si256(_mm256_srli_epi16(r0, 4), _mm256_set1_epi8(3)); + __m256i r3 = _mm256_and_si256(_mm256_srli_epi16(r0, 6), _mm256_set1_epi8(3)); + r0 = _mm256_and_si256(r0, _mm256_set1_epi8(3)); + b0 = _mm256_add_epi16(b0, _mm256_add_epi16(_mm256_maddubs_epi16(r0, qy0), _mm256_maddubs_epi16(r1, qy1))); + b1 = _mm256_add_epi16(b1, _mm256_add_epi16(_mm256_maddubs_epi16(r2, qy2), _mm256_maddubs_epi16(r3, qy3))); + + __m256i w0 = _mm256_loadu_si256((const __m256i *) (x[i].qs3 + j)); + __m256i w1 = _mm256_and_si256(_mm256_srli_epi16(w0, 2), _mm256_set1_epi8(3)); + __m256i w2 = _mm256_and_si256(_mm256_srli_epi16(w0, 4), _mm256_set1_epi8(3)); + __m256i w3 = _mm256_and_si256(_mm256_srli_epi16(w0, 6), _mm256_set1_epi8(3)); + w0 = _mm256_and_si256(w0, _mm256_set1_epi8(3)); + c0 = _mm256_add_epi16(c0, _mm256_add_epi16(_mm256_maddubs_epi16(w0, qy0), _mm256_maddubs_epi16(w1, qy1))); + c1 = _mm256_add_epi16(c1, _mm256_add_epi16(_mm256_maddubs_epi16(w2, qy2), _mm256_maddubs_epi16(w3, qy3))); + + __m256i v0 = _mm256_loadu_si256((const __m256i *) (x[i].qs4 + j)); + __m256i v1 = _mm256_and_si256(_mm256_srli_epi16(v0, 2), _mm256_set1_epi8(3)); + __m256i v2 = _mm256_and_si256(_mm256_srli_epi16(v0, 4), _mm256_set1_epi8(3)); + __m256i v3 = _mm256_and_si256(_mm256_srli_epi16(v0, 6), _mm256_set1_epi8(3)); + v0 = _mm256_and_si256(v0, _mm256_set1_epi8(3)); + e0 = _mm256_add_epi16(e0, _mm256_add_epi16(_mm256_maddubs_epi16(v0, qy0), _mm256_maddubs_epi16(v1, qy1))); + e1 = _mm256_add_epi16(e1, _mm256_add_epi16(_mm256_maddubs_epi16(v2, qy2), _mm256_maddubs_epi16(v3, qy3))); + } + + const __m256i ysum = _mm256_loadu_si256((const __m256i *) y[i].bsums); + const __m256 d1 = _mm256_set1_ps(y[i].d * GGML_CPU_FP16_TO_FP32(x[i].d1)); + const __m256 d2 = _mm256_set1_ps(y[i].d * GGML_CPU_FP16_TO_FP32(x[i].d2)); + const __m256 d3 = _mm256_set1_ps(y[i].d * GGML_CPU_FP16_TO_FP32(x[i].d3)); + const __m256 d4 = _mm256_set1_ps(y[i].d * GGML_CPU_FP16_TO_FP32(x[i].d4)); + + __m256i sa = _mm256_madd_epi16(_mm256_sub_epi16(_mm256_add_epi16(a0, a1), ysum), _mm256_set1_epi16(1)); + __m256i sb = _mm256_madd_epi16(_mm256_sub_epi16(_mm256_add_epi16(b0, b1), ysum), _mm256_set1_epi16(1)); + __m256i sc = _mm256_madd_epi16(_mm256_sub_epi16(_mm256_add_epi16(c0, c1), ysum), _mm256_set1_epi16(1)); + __m256i se = _mm256_madd_epi16(_mm256_sub_epi16(_mm256_add_epi16(e0, e1), ysum), _mm256_set1_epi16(1)); + + sumf = _mm256_add_ps(_mm256_mul_ps(_mm256_cvtepi32_ps(sa), d1), sumf); + sumf = _mm256_add_ps(_mm256_mul_ps(_mm256_cvtepi32_ps(sb), d2), sumf); + sumf = _mm256_add_ps(_mm256_mul_ps(_mm256_cvtepi32_ps(sc), d3), sumf); + sumf = _mm256_add_ps(_mm256_mul_ps(_mm256_cvtepi32_ps(se), d4), sumf); + } + + *s = hsum_float_8(sumf); +#else + UNUSED(x); + UNUSED(y); + UNUSED(nb); + ggml_vec_dot_tq4p_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} + +#if defined(__AVX2__) +static void tqkp_gv_vec_dot_avx2(int n, float * GGML_RESTRICT s, const void * GGML_RESTRICT vx, const void * GGML_RESTRICT vy, int K, int g) { + const int nsub = QK_TQKP_SUPER / g, qb = QK_TQKP_SUPER / 4, bpb = K*qb + K*nsub*2; + const int cps = g / 32; // chunks (of 32 weights) per scale sub-group: g=32/64/128 -> 1/2/4 + const block_q8_K * GGML_RESTRICT y = (const block_q8_K *) vy; + const uint8_t * GGML_RESTRICT x = (const uint8_t *) vx; + const int nb = n / QK_TQKP_SUPER; + const __m256i mask = _mm256_set1_epi8(3), ones = _mm256_set1_epi16(1); + __m256 sumf = _mm256_setzero_ps(); // partials stay in lanes; one hsum at the end (as the g=256 kernel) + for (int i = 0; i < nb; ++i) { + const uint8_t * GGML_RESTRICT blk = x + (size_t) i * bpb; + const ggml_half * GGML_RESTRICT d = (const ggml_half *)(blk + K*qb); + __m256i qy[8]; + for (int c = 0; c < 8; ++c) qy[c] = _mm256_loadu_si256((const __m256i *)(y[i].qs + c*32)); + const int16_t * GGML_RESTRICT bs = y[i].bsums; + for (int k = 0; k < K; ++k) { + const uint8_t * GGML_RESTRICT qs = blk + (size_t) k * qb; + __m256i m[8]; + __m256i pa = _mm256_loadu_si256((const __m256i *)(qs + 0)); + m[0] = _mm256_maddubs_epi16(_mm256_and_si256(pa, mask), qy[0]); + m[1] = _mm256_maddubs_epi16(_mm256_and_si256(_mm256_srli_epi16(pa, 2), mask), qy[1]); + m[2] = _mm256_maddubs_epi16(_mm256_and_si256(_mm256_srli_epi16(pa, 4), mask), qy[2]); + m[3] = _mm256_maddubs_epi16(_mm256_and_si256(_mm256_srli_epi16(pa, 6), mask), qy[3]); + __m256i pb = _mm256_loadu_si256((const __m256i *)(qs + 32)); + m[4] = _mm256_maddubs_epi16(_mm256_and_si256(pb, mask), qy[4]); + m[5] = _mm256_maddubs_epi16(_mm256_and_si256(_mm256_srli_epi16(pb, 2), mask), qy[5]); + m[6] = _mm256_maddubs_epi16(_mm256_and_si256(_mm256_srli_epi16(pb, 4), mask), qy[6]); + m[7] = _mm256_maddubs_epi16(_mm256_and_si256(_mm256_srli_epi16(pb, 6), mask), qy[7]); + for (int sub = 0; sub < nsub; ++sub) { + __m256i acc = _mm256_setzero_si256(); + int32_t ysum = 0; + for (int cc = 0; cc < cps; ++cc) { + const int c = sub*cps + cc; + acc = _mm256_add_epi16(acc, m[c]); + ysum += (int32_t) bs[c*2] + (int32_t) bs[c*2 + 1]; + } + // sum(u*y) partials in lanes; subtract the sub-group ysum from one lane -> sum(t*y) + __m256i r = _mm256_sub_epi32(_mm256_madd_epi16(acc, ones), _mm256_set_epi32(0, 0, 0, 0, 0, 0, 0, ysum)); + const float sc = y[i].d * GGML_CPU_FP16_TO_FP32(d[k*nsub + sub]); + sumf = _mm256_add_ps(sumf, _mm256_mul_ps(_mm256_cvtepi32_ps(r), _mm256_set1_ps(sc))); + } + } + } + *s = hsum_float_8(sumf); +} + +// sasori q4k6 (g=32) AVX2 vec_dot: trit-dot IDENTICAL to tqkp_gv_vec_dot_avx2 (K planes, nsub=8, +// cps=1); ONLY the per-sub-group scale decode differs (q4k6: flip fp32 sign bit iff sub sign set, +// then * u6, over a per-plane fp16 super). Layout per 256-super-block: K*64 trit bytes, then K*nsub +// sub bytes (bit7 sign, bits5..0 u6), then K fp16 super-scales. Gated scalar==AVX2 by bench_q4k6_q8k.c. +static void q4k6_vec_dot_avx2(int n, float * GGML_RESTRICT s, const void * GGML_RESTRICT vx, const void * GGML_RESTRICT vy, int K) { + const int nsub = QK_TQKP_SUPER / Q4K6_G, qb = QK_TQKP_SUPER / 4; // nsub=8, qb=64 + const int bpb = K*qb + K*nsub + K*2; // trits + sub bytes + super f16 + const block_q8_K * GGML_RESTRICT y = (const block_q8_K *) vy; + const uint8_t * GGML_RESTRICT x = (const uint8_t *) vx; + const int nb = n / QK_TQKP_SUPER; + const __m256i mask = _mm256_set1_epi8(3), ones = _mm256_set1_epi16(1); + __m256 sumf = _mm256_setzero_ps(); + for (int i = 0; i < nb; ++i) { + const uint8_t * GGML_RESTRICT blk = x + (size_t) i * bpb; + const uint8_t * GGML_RESTRICT sub = blk + K*qb; // K*nsub sub bytes + const ggml_half * GGML_RESTRICT sup = (const ggml_half *)(blk + K*qb + K*nsub); // K super f16 + __m256i qy[8]; + for (int c = 0; c < 8; ++c) qy[c] = _mm256_loadu_si256((const __m256i *)(y[i].qs + c*32)); + const int16_t * GGML_RESTRICT bs = y[i].bsums; + for (int k = 0; k < K; ++k) { + const uint8_t * GGML_RESTRICT qs = blk + (size_t) k * qb; + const float superf = GGML_CPU_FP16_TO_FP32(sup[k]); + __m256i m[8]; + __m256i pa = _mm256_loadu_si256((const __m256i *)(qs + 0)); + m[0] = _mm256_maddubs_epi16(_mm256_and_si256(pa, mask), qy[0]); + m[1] = _mm256_maddubs_epi16(_mm256_and_si256(_mm256_srli_epi16(pa, 2), mask), qy[1]); + m[2] = _mm256_maddubs_epi16(_mm256_and_si256(_mm256_srli_epi16(pa, 4), mask), qy[2]); + m[3] = _mm256_maddubs_epi16(_mm256_and_si256(_mm256_srli_epi16(pa, 6), mask), qy[3]); + __m256i pb = _mm256_loadu_si256((const __m256i *)(qs + 32)); + m[4] = _mm256_maddubs_epi16(_mm256_and_si256(pb, mask), qy[4]); + m[5] = _mm256_maddubs_epi16(_mm256_and_si256(_mm256_srli_epi16(pb, 2), mask), qy[5]); + m[6] = _mm256_maddubs_epi16(_mm256_and_si256(_mm256_srli_epi16(pb, 4), mask), qy[6]); + m[7] = _mm256_maddubs_epi16(_mm256_and_si256(_mm256_srli_epi16(pb, 6), mask), qy[7]); + for (int c = 0; c < nsub; ++c) { // g=32 -> cps=1, one 32-chunk per sub + const uint8_t sb = sub[k*nsub + c]; + uint32_t bits; memcpy(&bits, &superf, 4); // q4k6 scale = flip_sign(super,sub.sign)*u6 + bits ^= ((uint32_t)(sb & 0x80u)) << 24; + float sc_f; memcpy(&sc_f, &bits, 4); + sc_f *= (float)(sb & 0x3F); + const int32_t ysum = (int32_t) bs[c*2] + (int32_t) bs[c*2 + 1]; + __m256i r = _mm256_sub_epi32(_mm256_madd_epi16(m[c], ones), _mm256_set_epi32(0, 0, 0, 0, 0, 0, 0, ysum)); + const float sc = y[i].d * sc_f; + sumf = _mm256_add_ps(sumf, _mm256_mul_ps(_mm256_cvtepi32_ps(r), _mm256_set1_ps(sc))); + } + } + } + *s = hsum_float_8(sumf); +} +#endif + +void ggml_vec_dot_tq1p_g32_q8_K(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) { + GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); +#if defined(__AVX2__) + tqkp_gv_vec_dot_avx2(n, s, vx, vy, 1, 32); +#else + ggml_vec_dot_tq1p_g32_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} +void ggml_vec_dot_tq1p_g64_q8_K(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) { + GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); +#if defined(__AVX2__) + tqkp_gv_vec_dot_avx2(n, s, vx, vy, 1, 64); +#else + ggml_vec_dot_tq1p_g64_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} +void ggml_vec_dot_tq1p_g128_q8_K(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) { + GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); +#if defined(__AVX2__) + tqkp_gv_vec_dot_avx2(n, s, vx, vy, 1, 128); +#else + ggml_vec_dot_tq1p_g128_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} +void ggml_vec_dot_tq2p_g32_q8_K(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) { + GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); +#if defined(__AVX2__) + tqkp_gv_vec_dot_avx2(n, s, vx, vy, 2, 32); +#else + ggml_vec_dot_tq2p_g32_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} +void ggml_vec_dot_tq2p_g64_q8_K(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) { + GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); +#if defined(__AVX2__) + tqkp_gv_vec_dot_avx2(n, s, vx, vy, 2, 64); +#else + ggml_vec_dot_tq2p_g64_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} +void ggml_vec_dot_tq2p_g128_q8_K(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) { + GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); +#if defined(__AVX2__) + tqkp_gv_vec_dot_avx2(n, s, vx, vy, 2, 128); +#else + ggml_vec_dot_tq2p_g128_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} +void ggml_vec_dot_tq3p_g32_q8_K(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) { + GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); +#if defined(__AVX2__) + tqkp_gv_vec_dot_avx2(n, s, vx, vy, 3, 32); +#else + ggml_vec_dot_tq3p_g32_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} +void ggml_vec_dot_tq3p_g64_q8_K(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) { + GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); +#if defined(__AVX2__) + tqkp_gv_vec_dot_avx2(n, s, vx, vy, 3, 64); +#else + ggml_vec_dot_tq3p_g64_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} +void ggml_vec_dot_tq3p_g128_q8_K(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) { + GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); +#if defined(__AVX2__) + tqkp_gv_vec_dot_avx2(n, s, vx, vy, 3, 128); +#else + ggml_vec_dot_tq3p_g128_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} +void ggml_vec_dot_tq4p_g32_q8_K(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) { + GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); +#if defined(__AVX2__) + tqkp_gv_vec_dot_avx2(n, s, vx, vy, 4, 32); +#else + ggml_vec_dot_tq4p_g32_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} +void ggml_vec_dot_tq4p_g64_q8_K(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) { + GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); +#if defined(__AVX2__) + tqkp_gv_vec_dot_avx2(n, s, vx, vy, 4, 64); +#else + ggml_vec_dot_tq4p_g64_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} +void ggml_vec_dot_tq4p_g128_q8_K(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) { + GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); +#if defined(__AVX2__) + tqkp_gv_vec_dot_avx2(n, s, vx, vy, 4, 128); +#else + ggml_vec_dot_tq4p_g128_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} +void ggml_vec_dot_tq1p_q4k6_q8_K(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) { + GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); +#if defined(__AVX2__) + q4k6_vec_dot_avx2(n, s, vx, vy, 1); +#else + ggml_vec_dot_tq1p_q4k6_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} +void ggml_vec_dot_tq2p_q4k6_q8_K(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) { + GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); +#if defined(__AVX2__) + q4k6_vec_dot_avx2(n, s, vx, vy, 2); +#else + ggml_vec_dot_tq2p_q4k6_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} +void ggml_vec_dot_tq3p_q4k6_q8_K(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) { + GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); +#if defined(__AVX2__) + q4k6_vec_dot_avx2(n, s, vx, vy, 3); +#else + ggml_vec_dot_tq3p_q4k6_q8_K_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} void ggml_vec_dot_q2_K_q8_K(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) { assert(nrc == 1); UNUSED(nrc); diff --git a/src/ggml-cpu/ggml-cpu.c b/src/ggml-cpu/ggml-cpu.c index eb8341c9..53af1417 100644 --- a/src/ggml-cpu/ggml-cpu.c +++ b/src/ggml-cpu/ggml-cpu.c @@ -400,6 +400,126 @@ static const struct ggml_type_traits_cpu type_traits_cpu[GGML_TYPE_COUNT] = { .vec_dot_type = GGML_TYPE_Q8_K, .nrows = 1, }, + [GGML_TYPE_TQ2P] = { + .from_float = quantize_row_tq2p, + .vec_dot = ggml_vec_dot_tq2p_q8_K, + .vec_dot_type = GGML_TYPE_Q8_K, + .nrows = 1, + }, + [GGML_TYPE_TQ3P] = { + .from_float = quantize_row_tq3p, + .vec_dot = ggml_vec_dot_tq3p_q8_K, + .vec_dot_type = GGML_TYPE_Q8_K, + .nrows = 1, + }, + [GGML_TYPE_TQ1P] = { + .from_float = quantize_row_tq1p, + .vec_dot = ggml_vec_dot_tq1p_q8_K, + .vec_dot_type = GGML_TYPE_Q8_K, + .nrows = 1, + }, + [GGML_TYPE_TQ4P] = { + .from_float = quantize_row_tq4p, + .vec_dot = ggml_vec_dot_tq4p_q8_K, + .vec_dot_type = GGML_TYPE_Q8_K, + .nrows = 1, + }, + [GGML_TYPE_TQ1P_G32] = { + .from_float = quantize_row_tq1p_g32, + .vec_dot = ggml_vec_dot_tq1p_g32_q8_K, + .vec_dot_type = GGML_TYPE_Q8_K, + .nrows = 1, + }, + [GGML_TYPE_TQ1P_G64] = { + .from_float = quantize_row_tq1p_g64, + .vec_dot = ggml_vec_dot_tq1p_g64_q8_K, + .vec_dot_type = GGML_TYPE_Q8_K, + .nrows = 1, + }, + [GGML_TYPE_TQ1P_G128] = { + .from_float = quantize_row_tq1p_g128, + .vec_dot = ggml_vec_dot_tq1p_g128_q8_K, + .vec_dot_type = GGML_TYPE_Q8_K, + .nrows = 1, + }, + [GGML_TYPE_TQ2P_G32] = { + .from_float = quantize_row_tq2p_g32, + .vec_dot = ggml_vec_dot_tq2p_g32_q8_K, + .vec_dot_type = GGML_TYPE_Q8_K, + .nrows = 1, + }, + [GGML_TYPE_TQ2P_G64] = { + .from_float = quantize_row_tq2p_g64, + .vec_dot = ggml_vec_dot_tq2p_g64_q8_K, + .vec_dot_type = GGML_TYPE_Q8_K, + .nrows = 1, + }, + [GGML_TYPE_TQ2P_G128] = { + .from_float = quantize_row_tq2p_g128, + .vec_dot = ggml_vec_dot_tq2p_g128_q8_K, + .vec_dot_type = GGML_TYPE_Q8_K, + .nrows = 1, + }, + [GGML_TYPE_TQ3P_G32] = { + .from_float = quantize_row_tq3p_g32, + .vec_dot = ggml_vec_dot_tq3p_g32_q8_K, + .vec_dot_type = GGML_TYPE_Q8_K, + .nrows = 1, + }, + [GGML_TYPE_TQ3P_G64] = { + .from_float = quantize_row_tq3p_g64, + .vec_dot = ggml_vec_dot_tq3p_g64_q8_K, + .vec_dot_type = GGML_TYPE_Q8_K, + .nrows = 1, + }, + [GGML_TYPE_TQ3P_G128] = { + .from_float = quantize_row_tq3p_g128, + .vec_dot = ggml_vec_dot_tq3p_g128_q8_K, + .vec_dot_type = GGML_TYPE_Q8_K, + .nrows = 1, + }, + [GGML_TYPE_TQ4P_G32] = { + .from_float = quantize_row_tq4p_g32, + .vec_dot = ggml_vec_dot_tq4p_g32_q8_K, + .vec_dot_type = GGML_TYPE_Q8_K, + .nrows = 1, + }, + [GGML_TYPE_TQ4P_G64] = { + .from_float = quantize_row_tq4p_g64, + .vec_dot = ggml_vec_dot_tq4p_g64_q8_K, + .vec_dot_type = GGML_TYPE_Q8_K, + .nrows = 1, + }, + [GGML_TYPE_TQ4P_G128] = { + .from_float = quantize_row_tq4p_g128, + .vec_dot = ggml_vec_dot_tq4p_g128_q8_K, + .vec_dot_type = GGML_TYPE_Q8_K, + .nrows = 1, + }, + [GGML_TYPE_TQ1P_Q4K6] = { + .from_float = quantize_row_tq1p_q4k6, + .vec_dot = ggml_vec_dot_tq1p_q4k6_q8_K, + .vec_dot_type = GGML_TYPE_Q8_K, + .nrows = 1, + }, + [GGML_TYPE_TQ2P_Q4K6] = { + .from_float = quantize_row_tq2p_q4k6, + .vec_dot = ggml_vec_dot_tq2p_q4k6_q8_K, + .vec_dot_type = GGML_TYPE_Q8_K, + .nrows = 1, + }, + [GGML_TYPE_TQ3P_Q4K6] = { + .from_float = quantize_row_tq3p_q4k6, + .vec_dot = ggml_vec_dot_tq3p_q4k6_q8_K, + .vec_dot_type = GGML_TYPE_Q8_K, + .nrows = 1, + }, + [GGML_TYPE_TQ1PA] = { + .from_float = quantize_row_tq1pa, + .vec_dot = ggml_vec_dot_tq1pa_q8_K, + .vec_dot_type = GGML_TYPE_Q8_K, + .nrows = 1, + }, [GGML_TYPE_I32] = { .from_float = (ggml_from_float_t) ggml_cpu_fp32_to_i32, }, diff --git a/src/ggml-cpu/quants.c b/src/ggml-cpu/quants.c index e5f9a408..621a8573 100644 --- a/src/ggml-cpu/quants.c +++ b/src/ggml-cpu/quants.c @@ -112,6 +112,267 @@ void quantize_row_tq2_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, quantize_row_tq2_0_ref(x, y, k); } +// TQ2P weights come pre-quantized from sasori (Python); the C weight-quantizer is a loud stub. +void quantize_row_tq2p(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) { + GGML_UNUSED(x); GGML_UNUSED(vy); GGML_UNUSED(k); + GGML_ABORT("TQ2P quantization is done in Python (sasori), not in llama.cpp"); +} + +// Multiplication-free dot of a TQ2P weight row with an F32 activation (vec_dot_type = F32). +// Trits {-1,0,+1} -> add/sub of the activation; only 2 multiplies per block (the scales). +void ggml_vec_dot_tq2p_q8_K_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) { + assert(n % QK_TQ2P == 0); + GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); + const block_tq2p * GGML_RESTRICT x = vx; + const block_q8_K * GGML_RESTRICT y = vy; + const int nb = n / QK_TQ2P; // QK_TQ2P == QK_K -> one tq2p block per q8_K block + float sumf = 0.0f; + for (int i = 0; i < nb; ++i) { + // u = trit + 1 in {0,1,2}; sum(u*yq) - sum(yq) = sum(trit*yq). yq are int8 activations. + int32_t su1 = 0, su2 = 0, ysum = 0; + for (int j = 0; j < QK_TQ2P; ++j) { + const int byte = (j / 128) * 32 + (j % 32); // tq2_0 strided layout (per plane) + const int bp = (j / 32) % 4; + const int u1 = (x[i].qs1[byte] >> (bp * 2)) & 3; + const int u2 = (x[i].qs2[byte] >> (bp * 2)) & 3; + const int yq = y[i].qs[j]; + su1 += u1 * yq; + su2 += u2 * yq; + ysum += yq; + } + const float d = y[i].d; + sumf += d * (GGML_FP16_TO_FP32(x[i].d1) * (float) (su1 - ysum) + + GGML_FP16_TO_FP32(x[i].d2) * (float) (su2 - ysum)); + } + *s = sumf; +} + +// TQ3P weights come pre-quantized from sasori (Python); the C weight-quantizer is a loud stub. +void quantize_row_tq3p(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) { + GGML_UNUSED(x); GGML_UNUSED(vy); GGML_UNUSED(k); + GGML_ABORT("TQ3P quantization is done in Python (sasori), not in llama.cpp"); +} + +// Multiplication-free dot of a TQ3P weight row with a Q8_K activation block (= TQ2P + one plane). +void ggml_vec_dot_tq3p_q8_K_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) { + assert(n % QK_TQ3P == 0); + GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); + const block_tq3p * GGML_RESTRICT x = vx; + const block_q8_K * GGML_RESTRICT y = vy; + const int nb = n / QK_TQ3P; // QK_TQ3P == QK_K -> one tq3p block per q8_K block + float sumf = 0.0f; + for (int i = 0; i < nb; ++i) { + // u = trit + 1 in {0,1,2}; sum(u*yq) - sum(yq) = sum(trit*yq). yq are int8 activations. + int32_t su1 = 0, su2 = 0, su3 = 0, ysum = 0; + for (int j = 0; j < QK_TQ3P; ++j) { + const int byte = (j / 128) * 32 + (j % 32); // tq2_0 strided layout (per plane) + const int bp = (j / 32) % 4; + const int u1 = (x[i].qs1[byte] >> (bp * 2)) & 3; + const int u2 = (x[i].qs2[byte] >> (bp * 2)) & 3; + const int u3 = (x[i].qs3[byte] >> (bp * 2)) & 3; + const int yq = y[i].qs[j]; + su1 += u1 * yq; + su2 += u2 * yq; + su3 += u3 * yq; + ysum += yq; + } + const float d = y[i].d; + sumf += d * (GGML_FP16_TO_FP32(x[i].d1) * (float) (su1 - ysum) + + GGML_FP16_TO_FP32(x[i].d2) * (float) (su2 - ysum) + + GGML_FP16_TO_FP32(x[i].d3) * (float) (su3 - ysum)); + } + *s = sumf; +} + +void quantize_row_tq1p(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) { + GGML_UNUSED(x); GGML_UNUSED(vy); GGML_UNUSED(k); + GGML_ABORT("TQ1P quantization is done in Python (sasori), not in llama.cpp"); +} + +void ggml_vec_dot_tq1p_q8_K_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) { + assert(n % QK_TQ1P == 0); + GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); + const block_tq1p * GGML_RESTRICT x = vx; + const block_q8_K * GGML_RESTRICT y = vy; + const int nb = n / QK_TQ1P; + float sumf = 0.0f; + for (int i = 0; i < nb; ++i) { + int32_t su1 = 0, ysum = 0; + for (int j = 0; j < QK_TQ1P; ++j) { + const int byte = (j / 128) * 32 + (j % 32); + const int bp = (j / 32) % 4; + const int u1 = (x[i].qs1[byte] >> (bp * 2)) & 3; + const int yq = y[i].qs[j]; + su1 += u1 * yq; + ysum += yq; + } + const float d = y[i].d; + sumf += d * (GGML_FP16_TO_FP32(x[i].d) * (float) (su1 - ysum)); + } + *s = sumf; +} + +void quantize_row_tq1pa(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) { + GGML_UNUSED(x); GGML_UNUSED(vy); GGML_UNUSED(k); + GGML_ABORT("TQ1PA quantization is done in Python (sasori), not in llama.cpp"); +} + +void ggml_vec_dot_tq1pa_q8_K_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) { + assert(n % QK_TQ1P == 0); + GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); + const block_tq1pa * GGML_RESTRICT x = vx; + const block_q8_K * GGML_RESTRICT y = vy; + const int nb = n / QK_TQ1P; + float sumf = 0.0f; + for (int i = 0; i < nb; ++i) { + int32_t su1 = 0, ysum = 0; + for (int j = 0; j < QK_TQ1P; ++j) { + const int byte = (j / 128) * 32 + (j % 32); + const int bp = (j / 32) % 4; + const int u1 = (x[i].qs1[byte] >> (bp * 2)) & 3; + const int yq = y[i].qs[j]; + su1 += u1 * yq; + ysum += yq; + } + const float d = y[i].d; + // asym: wh = d1*trit + mu => dot = d8*( d1*(su1-ysum) + mu*ysum ). g256: ysum-grupo == ysum-bloco. + sumf += d * (GGML_FP16_TO_FP32(x[i].d) * (float) (su1 - ysum) + + GGML_FP16_TO_FP32(x[i].mu) * (float) ysum); + } + *s = sumf; +} + +void quantize_row_tq4p(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) { + GGML_UNUSED(x); GGML_UNUSED(vy); GGML_UNUSED(k); + GGML_ABORT("TQ4P quantization is done in Python (sasori), not in llama.cpp"); +} + +void ggml_vec_dot_tq4p_q8_K_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) { + assert(n % QK_TQ4P == 0); + GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); + const block_tq4p * GGML_RESTRICT x = vx; + const block_q8_K * GGML_RESTRICT y = vy; + const int nb = n / QK_TQ4P; + float sumf = 0.0f; + for (int i = 0; i < nb; ++i) { + int32_t su1 = 0, su2 = 0, su3 = 0, su4 = 0, ysum = 0; + for (int j = 0; j < QK_TQ4P; ++j) { + const int byte = (j / 128) * 32 + (j % 32); + const int bp = (j / 32) % 4; + const int u1 = (x[i].qs1[byte] >> (bp * 2)) & 3; + const int u2 = (x[i].qs2[byte] >> (bp * 2)) & 3; + const int u3 = (x[i].qs3[byte] >> (bp * 2)) & 3; + const int u4 = (x[i].qs4[byte] >> (bp * 2)) & 3; + const int yq = y[i].qs[j]; + su1 += u1 * yq; + su2 += u2 * yq; + su3 += u3 * yq; + su4 += u4 * yq; + ysum += yq; + } + const float d = y[i].d; + sumf += d * (GGML_FP16_TO_FP32(x[i].d1) * (float) (su1 - ysum) + + GGML_FP16_TO_FP32(x[i].d2) * (float) (su2 - ysum) + + GGML_FP16_TO_FP32(x[i].d3) * (float) (su3 - ysum) + + GGML_FP16_TO_FP32(x[i].d4) * (float) (su4 - ysum)); + } + *s = sumf; +} + +// sasori g-variable vec_dot, scalar: trit*activation, sub-scale d[k][p/g] per group. AVX2 in arch/x86. +static void tqkp_gv_vec_dot(int n, float * GGML_RESTRICT s, const void * GGML_RESTRICT vx, const void * GGML_RESTRICT vy, int K, int g) { + assert(n % QK_TQKP_SUPER == 0); + const int nsub = QK_TQKP_SUPER / g, qb = QK_TQKP_SUPER / 4, bpb = K*qb + K*nsub*2; + const block_q8_K * GGML_RESTRICT y = (const block_q8_K *) vy; + const uint8_t * GGML_RESTRICT x = (const uint8_t *) vx; + const int nb = n / QK_TQKP_SUPER; + float sumf = 0.0f; + for (int i = 0; i < nb; ++i) { + const uint8_t * GGML_RESTRICT blk = x + (size_t) i * bpb; + const ggml_half * GGML_RESTRICT d = (const ggml_half *)(blk + K*qb); + for (int kk = 0; kk < K; ++kk) { + for (int sub = 0; sub < nsub; ++sub) { + int32_t acc = 0; + for (int pp = 0; pp < g; ++pp) { + const int p = sub*g + pp; + const int byte = (p/128)*32 + (p%32), bp = (p/32)%4; + const int t = ((blk[kk*qb + byte] >> (bp*2)) & 3) - 1; + acc += t * (int) y[i].qs[p]; + } + sumf += y[i].d * GGML_FP16_TO_FP32(d[kk*nsub + sub]) * (float) acc; + } + } + } + *s = sumf; +} + +void ggml_vec_dot_tq1p_g32_q8_K_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) { GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); tqkp_gv_vec_dot(n, s, vx, vy, 1, 32); } +void quantize_row_tq1p_g32(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) { GGML_UNUSED(x); GGML_UNUSED(vy); GGML_UNUSED(k); GGML_ABORT("TQ1P_G32 quantization is done in Python (sasori)"); } +void ggml_vec_dot_tq1p_g64_q8_K_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) { GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); tqkp_gv_vec_dot(n, s, vx, vy, 1, 64); } +void quantize_row_tq1p_g64(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) { GGML_UNUSED(x); GGML_UNUSED(vy); GGML_UNUSED(k); GGML_ABORT("TQ1P_G64 quantization is done in Python (sasori)"); } +void ggml_vec_dot_tq1p_g128_q8_K_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) { GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); tqkp_gv_vec_dot(n, s, vx, vy, 1, 128); } +void quantize_row_tq1p_g128(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) { GGML_UNUSED(x); GGML_UNUSED(vy); GGML_UNUSED(k); GGML_ABORT("TQ1P_G128 quantization is done in Python (sasori)"); } +void ggml_vec_dot_tq2p_g32_q8_K_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) { GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); tqkp_gv_vec_dot(n, s, vx, vy, 2, 32); } +void quantize_row_tq2p_g32(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) { GGML_UNUSED(x); GGML_UNUSED(vy); GGML_UNUSED(k); GGML_ABORT("TQ2P_G32 quantization is done in Python (sasori)"); } +void ggml_vec_dot_tq2p_g64_q8_K_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) { GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); tqkp_gv_vec_dot(n, s, vx, vy, 2, 64); } +void quantize_row_tq2p_g64(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) { GGML_UNUSED(x); GGML_UNUSED(vy); GGML_UNUSED(k); GGML_ABORT("TQ2P_G64 quantization is done in Python (sasori)"); } +void ggml_vec_dot_tq2p_g128_q8_K_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) { GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); tqkp_gv_vec_dot(n, s, vx, vy, 2, 128); } +void quantize_row_tq2p_g128(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) { GGML_UNUSED(x); GGML_UNUSED(vy); GGML_UNUSED(k); GGML_ABORT("TQ2P_G128 quantization is done in Python (sasori)"); } +void ggml_vec_dot_tq3p_g32_q8_K_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) { GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); tqkp_gv_vec_dot(n, s, vx, vy, 3, 32); } +void quantize_row_tq3p_g32(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) { GGML_UNUSED(x); GGML_UNUSED(vy); GGML_UNUSED(k); GGML_ABORT("TQ3P_G32 quantization is done in Python (sasori)"); } +void ggml_vec_dot_tq3p_g64_q8_K_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) { GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); tqkp_gv_vec_dot(n, s, vx, vy, 3, 64); } +void quantize_row_tq3p_g64(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) { GGML_UNUSED(x); GGML_UNUSED(vy); GGML_UNUSED(k); GGML_ABORT("TQ3P_G64 quantization is done in Python (sasori)"); } +void ggml_vec_dot_tq3p_g128_q8_K_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) { GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); tqkp_gv_vec_dot(n, s, vx, vy, 3, 128); } +void quantize_row_tq3p_g128(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) { GGML_UNUSED(x); GGML_UNUSED(vy); GGML_UNUSED(k); GGML_ABORT("TQ3P_G128 quantization is done in Python (sasori)"); } +void ggml_vec_dot_tq4p_g32_q8_K_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) { GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); tqkp_gv_vec_dot(n, s, vx, vy, 4, 32); } +void quantize_row_tq4p_g32(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) { GGML_UNUSED(x); GGML_UNUSED(vy); GGML_UNUSED(k); GGML_ABORT("TQ4P_G32 quantization is done in Python (sasori)"); } +void ggml_vec_dot_tq4p_g64_q8_K_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) { GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); tqkp_gv_vec_dot(n, s, vx, vy, 4, 64); } +void quantize_row_tq4p_g64(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) { GGML_UNUSED(x); GGML_UNUSED(vy); GGML_UNUSED(k); GGML_ABORT("TQ4P_G64 quantization is done in Python (sasori)"); } +void ggml_vec_dot_tq4p_g128_q8_K_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) { GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); tqkp_gv_vec_dot(n, s, vx, vy, 4, 128); } +void quantize_row_tq4p_g128(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) { GGML_UNUSED(x); GGML_UNUSED(vy); GGML_UNUSED(k); GGML_ABORT("TQ4P_G128 quantization is done in Python (sasori)"); } + +// sasori q4k6 (g=32) scalar vec_dot: trit-dot IDENTICAL to tqkp_gv_vec_dot (256-strided planes), ONLY +// the per-sub-group scale decode differs: q4k6 = flip fp32 sign bit iff sub sign set, then * u6, over a +// per-plane fp16 super. Layout/256: K*64 trit bytes, K*nsub sub bytes (bit7 sign, bits5..0 u6), K super f16. +static void q4k6_vec_dot(int n, float * GGML_RESTRICT s, const void * GGML_RESTRICT vx, const void * GGML_RESTRICT vy, int K) { + assert(n % QK_TQKP_SUPER == 0); + const int nsub = QK_TQKP_SUPER / Q4K6_G, qb = QK_TQKP_SUPER / 4, bpb = K*qb + K*nsub + K*2; + const block_q8_K * GGML_RESTRICT y = (const block_q8_K *) vy; + const uint8_t * GGML_RESTRICT x = (const uint8_t *) vx; + const int nb = n / QK_TQKP_SUPER; + float sumf = 0.0f; + for (int i = 0; i < nb; ++i) { + const uint8_t * GGML_RESTRICT blk = x + (size_t) i * bpb; + const uint8_t * GGML_RESTRICT sub = blk + K*qb; // K*nsub sub bytes + const ggml_half * GGML_RESTRICT sup = (const ggml_half *)(blk + K*qb + K*nsub); // K super f16 + for (int kk = 0; kk < K; ++kk) { + const float superf = GGML_FP16_TO_FP32(sup[kk]); + for (int c = 0; c < nsub; ++c) { + const uint8_t sb = sub[kk*nsub + c]; + uint32_t bits; memcpy(&bits, &superf, 4); // q4k6 scale = flip_sign(super,sb.sign)*u6 + bits ^= ((uint32_t)(sb & 0x80u)) << 24; + float sc; memcpy(&sc, &bits, 4); + sc *= (float)(sb & 0x3F); + int32_t acc = 0; + for (int pp = 0; pp < Q4K6_G; ++pp) { + const int p = c*Q4K6_G + pp; + const int byte = (p/128)*32 + (p%32), bp = (p/32)%4; + const int t = ((blk[kk*qb + byte] >> (bp*2)) & 3) - 1; + acc += t * (int) y[i].qs[p]; + } + sumf += y[i].d * sc * (float) acc; + } + } + } + *s = sumf; +} + +void ggml_vec_dot_tq1p_q4k6_q8_K_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) { GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); q4k6_vec_dot(n, s, vx, vy, 1); } +void quantize_row_tq1p_q4k6(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) { GGML_UNUSED(x); GGML_UNUSED(vy); GGML_UNUSED(k); GGML_ABORT("TQ1P_Q4K6 quantization is done in Python (sasori)"); } +void ggml_vec_dot_tq2p_q4k6_q8_K_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) { GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); q4k6_vec_dot(n, s, vx, vy, 2); } +void quantize_row_tq2p_q4k6(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) { GGML_UNUSED(x); GGML_UNUSED(vy); GGML_UNUSED(k); GGML_ABORT("TQ2P_Q4K6 quantization is done in Python (sasori)"); } +void ggml_vec_dot_tq3p_q4k6_q8_K_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) { GGML_UNUSED(bs); GGML_UNUSED(bx); GGML_UNUSED(by); GGML_UNUSED(nrc); q4k6_vec_dot(n, s, vx, vy, 3); } +void quantize_row_tq3p_q4k6(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) { GGML_UNUSED(x); GGML_UNUSED(vy); GGML_UNUSED(k); GGML_ABORT("TQ3P_Q4K6 quantization is done in Python (sasori)"); } //===================================== Q8_K ============================================== void quantize_row_q8_K_generic(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { diff --git a/src/ggml-cpu/quants.h b/src/ggml-cpu/quants.h index d4bc87a1..304127a5 100644 --- a/src/ggml-cpu/quants.h +++ b/src/ggml-cpu/quants.h @@ -32,6 +32,26 @@ void quantize_row_q8_K(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, in void quantize_row_tq1_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_tq2_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_tq2p(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_tq3p(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_tq1p(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_tq1pa(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_tq4p(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_tq1p_g32(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_tq1p_g64(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_tq1p_g128(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_tq2p_g32(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_tq2p_g64(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_tq2p_g128(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_tq3p_g32(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_tq3p_g64(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_tq3p_g128(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_tq4p_g32(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_tq4p_g64(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_tq4p_g128(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_tq1p_q4k6(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_tq2p_q4k6(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_tq3p_q4k6(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_iq4_nl (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_iq4_xs (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); @@ -55,6 +75,26 @@ void ggml_vec_dot_q6_K_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const voi void ggml_vec_dot_tq1_0_q8_K(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_tq2_0_q8_K(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_tq2p_q8_K(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_tq3p_q8_K(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_tq1p_q8_K(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_tq1pa_q8_K(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_tq4p_q8_K(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_tq1p_g32_q8_K(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_tq1p_g64_q8_K(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_tq1p_g128_q8_K(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_tq2p_g32_q8_K(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_tq2p_g64_q8_K(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_tq2p_g128_q8_K(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_tq3p_g32_q8_K(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_tq3p_g64_q8_K(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_tq3p_g128_q8_K(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_tq4p_g32_q8_K(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_tq4p_g64_q8_K(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_tq4p_g128_q8_K(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_tq1p_q4k6_q8_K(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_tq2p_q4k6_q8_K(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_tq3p_q4k6_q8_K(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_iq2_xxs_q8_K(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_iq2_xs_q8_K (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); @@ -82,6 +122,26 @@ void ggml_vec_dot_nvfp4_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, void ggml_vec_dot_tq1_0_q8_K_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_tq2_0_q8_K_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_tq2p_q8_K_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_tq3p_q8_K_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_tq1p_q8_K_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_tq1pa_q8_K_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_tq4p_q8_K_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_tq1p_g32_q8_K_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_tq1p_g64_q8_K_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_tq1p_g128_q8_K_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_tq2p_g32_q8_K_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_tq2p_g64_q8_K_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_tq2p_g128_q8_K_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_tq3p_g32_q8_K_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_tq3p_g64_q8_K_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_tq3p_g128_q8_K_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_tq4p_g32_q8_K_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_tq4p_g64_q8_K_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_tq4p_g128_q8_K_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_tq1p_q4k6_q8_K_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_tq2p_q4k6_q8_K_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_tq3p_q4k6_q8_K_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_K_q8_K_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_q3_K_q8_K_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/src/ggml-quants.c b/src/ggml-quants.c index 15d231f7..67477eda 100644 --- a/src/ggml-quants.c +++ b/src/ggml-quants.c @@ -2411,6 +2411,233 @@ void dequantize_row_tq2_0(const block_tq2_0 * GGML_RESTRICT x, float * GGML_REST } } +// sasori TQ2P dual-trit-plane dequant (layout = sasori/pack.py: 4 trits/byte little-endian per plane) +void dequantize_row_tq2p(const block_tq2p * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + assert(k % QK_TQ2P == 0); + const int64_t nb = k / QK_TQ2P; + for (int64_t i = 0; i < nb; ++i) { + const float d1 = GGML_FP16_TO_FP32(x[i].d1); + const float d2 = GGML_FP16_TO_FP32(x[i].d2); + // tq2_0 strided layout (per plane): logical position j -> byte (j/128)*32 + (j%32), + // bit-pair (j/32)%4. Must match sasori pack.py `_pack_plane_strided`. + for (int j = 0; j < QK_TQ2P; ++j) { + const int byte = (j / 128) * 32 + (j % 32); + const int bp = (j / 32) % 4; + const int t1 = ((x[i].qs1[byte] >> (bp * 2)) & 3) - 1; + const int t2 = ((x[i].qs2[byte] >> (bp * 2)) & 3) - 1; + y[i*QK_TQ2P + j] = d1 * (float) t1 + d2 * (float) t2; + } + } +} + +// TQ2P weights are produced by sasori (Python joint-ridge), never by llama.cpp -> quantize is a loud stub. +void quantize_row_tq2p_ref(const float * GGML_RESTRICT x, block_tq2p * GGML_RESTRICT y, int64_t k) { + (void) x; (void) y; (void) k; + GGML_ABORT("TQ2P quantization is done in Python (sasori), not in llama.cpp"); +} +size_t quantize_tq2p(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { + (void) src; (void) dst; (void) nrow; (void) n_per_row; (void) quant_weights; + GGML_ABORT("TQ2P quantization is done in Python (sasori), not in llama.cpp"); +} + +// sasori TQ3P tri-trit-plane dequant (= block_tq2p + one extra plane; same strided layout per plane) +void dequantize_row_tq3p(const block_tq3p * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + assert(k % QK_TQ3P == 0); + const int64_t nb = k / QK_TQ3P; + for (int64_t i = 0; i < nb; ++i) { + const float d1 = GGML_FP16_TO_FP32(x[i].d1); + const float d2 = GGML_FP16_TO_FP32(x[i].d2); + const float d3 = GGML_FP16_TO_FP32(x[i].d3); + // tq2_0 strided layout (per plane): logical position j -> byte (j/128)*32 + (j%32), + // bit-pair (j/32)%4. Must match sasori pack.py `_pack_plane_strided`. + for (int j = 0; j < QK_TQ3P; ++j) { + const int byte = (j / 128) * 32 + (j % 32); + const int bp = (j / 32) % 4; + const int t1 = ((x[i].qs1[byte] >> (bp * 2)) & 3) - 1; + const int t2 = ((x[i].qs2[byte] >> (bp * 2)) & 3) - 1; + const int t3 = ((x[i].qs3[byte] >> (bp * 2)) & 3) - 1; + y[i*QK_TQ3P + j] = d1 * (float) t1 + d2 * (float) t2 + d3 * (float) t3; + } + } +} + +// TQ3P weights are produced by sasori (Python joint-ridge), never by llama.cpp -> quantize is a loud stub. +void quantize_row_tq3p_ref(const float * GGML_RESTRICT x, block_tq3p * GGML_RESTRICT y, int64_t k) { + (void) x; (void) y; (void) k; + GGML_ABORT("TQ3P quantization is done in Python (sasori), not in llama.cpp"); +} +size_t quantize_tq3p(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { + (void) src; (void) dst; (void) nrow; (void) n_per_row; (void) quant_weights; + GGML_ABORT("TQ3P quantization is done in Python (sasori), not in llama.cpp"); +} + +void dequantize_row_tq1p(const block_tq1p * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + assert(k % QK_TQ1P == 0); + const int64_t nb = k / QK_TQ1P; + for (int64_t i = 0; i < nb; ++i) { + const float d = GGML_FP16_TO_FP32(x[i].d); + for (int j = 0; j < QK_TQ1P; ++j) { + const int byte = (j / 128) * 32 + (j % 32); + const int bp = (j / 32) % 4; + const int t1 = ((x[i].qs1[byte] >> (bp * 2)) & 3) - 1; + y[i*QK_TQ1P + j] = d * (float) t1; + } + } +} + +void quantize_row_tq1p_ref(const float * GGML_RESTRICT x, block_tq1p * GGML_RESTRICT y, int64_t k) { + (void) x; (void) y; (void) k; + GGML_ABORT("TQ1P quantization is done in Python (sasori), not in llama.cpp"); +} +size_t quantize_tq1p(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { + (void) src; (void) dst; (void) nrow; (void) n_per_row; (void) quant_weights; + GGML_ABORT("TQ1P quantization is done in Python (sasori), not in llama.cpp"); +} + +// sasori TQ1PA = block_tq1p + offset μ f16 do grupo (g256). Decode: wh = d*trit + μ (soma μ a todo peso). +void dequantize_row_tq1pa(const block_tq1pa * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + assert(k % QK_TQ1P == 0); + const int64_t nb = k / QK_TQ1P; + for (int64_t i = 0; i < nb; ++i) { + const float d = GGML_FP16_TO_FP32(x[i].d); + const float mu = GGML_FP16_TO_FP32(x[i].mu); + for (int j = 0; j < QK_TQ1P; ++j) { + const int byte = (j / 128) * 32 + (j % 32); + const int bp = (j / 32) % 4; + const int t1 = ((x[i].qs1[byte] >> (bp * 2)) & 3) - 1; + y[i*QK_TQ1P + j] = d * (float) t1 + mu; + } + } +} + +void quantize_row_tq1pa_ref(const float * GGML_RESTRICT x, block_tq1pa * GGML_RESTRICT y, int64_t k) { + (void) x; (void) y; (void) k; + GGML_ABORT("TQ1PA quantization is done in Python (sasori), not in llama.cpp"); +} + +void dequantize_row_tq4p(const block_tq4p * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + assert(k % QK_TQ4P == 0); + const int64_t nb = k / QK_TQ4P; + for (int64_t i = 0; i < nb; ++i) { + const float d1 = GGML_FP16_TO_FP32(x[i].d1); + const float d2 = GGML_FP16_TO_FP32(x[i].d2); + const float d3 = GGML_FP16_TO_FP32(x[i].d3); + const float d4 = GGML_FP16_TO_FP32(x[i].d4); + for (int j = 0; j < QK_TQ4P; ++j) { + const int byte = (j / 128) * 32 + (j % 32); + const int bp = (j / 32) % 4; + const int t1 = ((x[i].qs1[byte] >> (bp * 2)) & 3) - 1; + const int t2 = ((x[i].qs2[byte] >> (bp * 2)) & 3) - 1; + const int t3 = ((x[i].qs3[byte] >> (bp * 2)) & 3) - 1; + const int t4 = ((x[i].qs4[byte] >> (bp * 2)) & 3) - 1; + y[i*QK_TQ4P + j] = d1 * (float) t1 + d2 * (float) t2 + d3 * (float) t3 + d4 * (float) t4; + } + } +} + +void quantize_row_tq4p_ref(const float * GGML_RESTRICT x, block_tq4p * GGML_RESTRICT y, int64_t k) { + (void) x; (void) y; (void) k; + GGML_ABORT("TQ4P quantization is done in Python (sasori), not in llama.cpp"); +} +size_t quantize_tq4p(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { + (void) src; (void) dst; (void) nrow; (void) n_per_row; (void) quant_weights; + GGML_ABORT("TQ4P quantization is done in Python (sasori), not in llama.cpp"); +} + +// sasori g-variable TQ{K}P_G{g} dequant: super-block 256, sub-scale d[k][p/g]; strided planes. +static void tqkp_gv_dequantize(const uint8_t * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k, int K, int g) { + assert(k % QK_TQKP_SUPER == 0); + const int nsub = QK_TQKP_SUPER / g, qb = QK_TQKP_SUPER / 4, bpb = K*qb + K*nsub*2; + const int64_t nb = k / QK_TQKP_SUPER; + for (int64_t i = 0; i < nb; ++i) { + const uint8_t * GGML_RESTRICT blk = x + (size_t) i * bpb; + const ggml_half * GGML_RESTRICT d = (const ggml_half *)(blk + K*qb); + for (int p = 0; p < QK_TQKP_SUPER; ++p) { + const int byte = (p/128)*32 + (p%32), bp = (p/32)%4, sub = p/g; + float acc = 0.0f; + for (int kk = 0; kk < K; ++kk) { + const int t = ((blk[kk*qb + byte] >> (bp*2)) & 3) - 1; + acc += GGML_FP16_TO_FP32(d[kk*nsub + sub]) * (float) t; + } + y[i*QK_TQKP_SUPER + p] = acc; + } + } +} + +void dequantize_row_tq1p_g32(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { tqkp_gv_dequantize((const uint8_t *) x, y, k, 1, 32); } +void quantize_row_tq1p_g32_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { (void) x; (void) y; (void) k; GGML_ABORT("TQ1P_G32 quantization is done in Python (sasori)"); } +size_t quantize_tq1p_g32(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { (void) src; (void) dst; (void) nrow; (void) n_per_row; (void) quant_weights; GGML_ABORT("TQ1P_G32 quantization is done in Python (sasori)"); } +void dequantize_row_tq1p_g64(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { tqkp_gv_dequantize((const uint8_t *) x, y, k, 1, 64); } +void quantize_row_tq1p_g64_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { (void) x; (void) y; (void) k; GGML_ABORT("TQ1P_G64 quantization is done in Python (sasori)"); } +size_t quantize_tq1p_g64(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { (void) src; (void) dst; (void) nrow; (void) n_per_row; (void) quant_weights; GGML_ABORT("TQ1P_G64 quantization is done in Python (sasori)"); } +void dequantize_row_tq1p_g128(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { tqkp_gv_dequantize((const uint8_t *) x, y, k, 1, 128); } +void quantize_row_tq1p_g128_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { (void) x; (void) y; (void) k; GGML_ABORT("TQ1P_G128 quantization is done in Python (sasori)"); } +size_t quantize_tq1p_g128(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { (void) src; (void) dst; (void) nrow; (void) n_per_row; (void) quant_weights; GGML_ABORT("TQ1P_G128 quantization is done in Python (sasori)"); } +void dequantize_row_tq2p_g32(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { tqkp_gv_dequantize((const uint8_t *) x, y, k, 2, 32); } +void quantize_row_tq2p_g32_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { (void) x; (void) y; (void) k; GGML_ABORT("TQ2P_G32 quantization is done in Python (sasori)"); } +size_t quantize_tq2p_g32(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { (void) src; (void) dst; (void) nrow; (void) n_per_row; (void) quant_weights; GGML_ABORT("TQ2P_G32 quantization is done in Python (sasori)"); } +void dequantize_row_tq2p_g64(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { tqkp_gv_dequantize((const uint8_t *) x, y, k, 2, 64); } +void quantize_row_tq2p_g64_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { (void) x; (void) y; (void) k; GGML_ABORT("TQ2P_G64 quantization is done in Python (sasori)"); } +size_t quantize_tq2p_g64(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { (void) src; (void) dst; (void) nrow; (void) n_per_row; (void) quant_weights; GGML_ABORT("TQ2P_G64 quantization is done in Python (sasori)"); } +void dequantize_row_tq2p_g128(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { tqkp_gv_dequantize((const uint8_t *) x, y, k, 2, 128); } +void quantize_row_tq2p_g128_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { (void) x; (void) y; (void) k; GGML_ABORT("TQ2P_G128 quantization is done in Python (sasori)"); } +size_t quantize_tq2p_g128(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { (void) src; (void) dst; (void) nrow; (void) n_per_row; (void) quant_weights; GGML_ABORT("TQ2P_G128 quantization is done in Python (sasori)"); } +void dequantize_row_tq3p_g32(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { tqkp_gv_dequantize((const uint8_t *) x, y, k, 3, 32); } +void quantize_row_tq3p_g32_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { (void) x; (void) y; (void) k; GGML_ABORT("TQ3P_G32 quantization is done in Python (sasori)"); } +size_t quantize_tq3p_g32(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { (void) src; (void) dst; (void) nrow; (void) n_per_row; (void) quant_weights; GGML_ABORT("TQ3P_G32 quantization is done in Python (sasori)"); } +void dequantize_row_tq3p_g64(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { tqkp_gv_dequantize((const uint8_t *) x, y, k, 3, 64); } +void quantize_row_tq3p_g64_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { (void) x; (void) y; (void) k; GGML_ABORT("TQ3P_G64 quantization is done in Python (sasori)"); } +size_t quantize_tq3p_g64(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { (void) src; (void) dst; (void) nrow; (void) n_per_row; (void) quant_weights; GGML_ABORT("TQ3P_G64 quantization is done in Python (sasori)"); } +void dequantize_row_tq3p_g128(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { tqkp_gv_dequantize((const uint8_t *) x, y, k, 3, 128); } +void quantize_row_tq3p_g128_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { (void) x; (void) y; (void) k; GGML_ABORT("TQ3P_G128 quantization is done in Python (sasori)"); } +size_t quantize_tq3p_g128(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { (void) src; (void) dst; (void) nrow; (void) n_per_row; (void) quant_weights; GGML_ABORT("TQ3P_G128 quantization is done in Python (sasori)"); } +void dequantize_row_tq4p_g32(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { tqkp_gv_dequantize((const uint8_t *) x, y, k, 4, 32); } +void quantize_row_tq4p_g32_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { (void) x; (void) y; (void) k; GGML_ABORT("TQ4P_G32 quantization is done in Python (sasori)"); } +size_t quantize_tq4p_g32(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { (void) src; (void) dst; (void) nrow; (void) n_per_row; (void) quant_weights; GGML_ABORT("TQ4P_G32 quantization is done in Python (sasori)"); } +void dequantize_row_tq4p_g64(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { tqkp_gv_dequantize((const uint8_t *) x, y, k, 4, 64); } +void quantize_row_tq4p_g64_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { (void) x; (void) y; (void) k; GGML_ABORT("TQ4P_G64 quantization is done in Python (sasori)"); } +size_t quantize_tq4p_g64(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { (void) src; (void) dst; (void) nrow; (void) n_per_row; (void) quant_weights; GGML_ABORT("TQ4P_G64 quantization is done in Python (sasori)"); } +void dequantize_row_tq4p_g128(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { tqkp_gv_dequantize((const uint8_t *) x, y, k, 4, 128); } +void quantize_row_tq4p_g128_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { (void) x; (void) y; (void) k; GGML_ABORT("TQ4P_G128 quantization is done in Python (sasori)"); } +size_t quantize_tq4p_g128(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { (void) src; (void) dst; (void) nrow; (void) n_per_row; (void) quant_weights; GGML_ABORT("TQ4P_G128 quantization is done in Python (sasori)"); } + +// sasori q4k6 (g=32) dequant: clone of tqkp_gv_dequantize (256-strided planes); ONLY the scale decode +// differs (q4k6 = flip fp32 sign bit iff sub sign set, then * u6, over a per-plane fp16 super). +static void q4k6_dequantize(const uint8_t * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k, int K) { + assert(k % QK_TQKP_SUPER == 0); + const int nsub = QK_TQKP_SUPER / Q4K6_G, qb = QK_TQKP_SUPER / 4, bpb = K*qb + K*nsub + K*2; + const int64_t nb = k / QK_TQKP_SUPER; + for (int64_t i = 0; i < nb; ++i) { + const uint8_t * GGML_RESTRICT blk = x + (size_t) i * bpb; + const uint8_t * GGML_RESTRICT sub = blk + K*qb; // K*nsub sub bytes + const ggml_half * GGML_RESTRICT sup = (const ggml_half *)(blk + K*qb + K*nsub); // K super f16 + for (int p = 0; p < QK_TQKP_SUPER; ++p) { + const int byte = (p/128)*32 + (p%32), bp = (p/32)%4, c = p/Q4K6_G; + float acc = 0.0f; + for (int kk = 0; kk < K; ++kk) { + const uint8_t sb = sub[kk*nsub + c]; + float superf = GGML_FP16_TO_FP32(sup[kk]); + uint32_t bits; memcpy(&bits, &superf, 4); + bits ^= ((uint32_t)(sb & 0x80u)) << 24; + float sc; memcpy(&sc, &bits, 4); + sc *= (float)(sb & 0x3F); + const int t = ((blk[kk*qb + byte] >> (bp*2)) & 3) - 1; + acc += sc * (float) t; + } + y[i*QK_TQKP_SUPER + p] = acc; + } + } +} +void dequantize_row_tq1p_q4k6(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { q4k6_dequantize((const uint8_t *) x, y, k, 1); } +void quantize_row_tq1p_q4k6_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { (void) x; (void) y; (void) k; GGML_ABORT("TQ1P_Q4K6 quantization is done in Python (sasori)"); } +size_t quantize_tq1p_q4k6(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { (void) src; (void) dst; (void) nrow; (void) n_per_row; (void) quant_weights; GGML_ABORT("TQ1P_Q4K6 quantization is done in Python (sasori)"); } +void dequantize_row_tq2p_q4k6(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { q4k6_dequantize((const uint8_t *) x, y, k, 2); } +void quantize_row_tq2p_q4k6_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { (void) x; (void) y; (void) k; GGML_ABORT("TQ2P_Q4K6 quantization is done in Python (sasori)"); } +size_t quantize_tq2p_q4k6(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { (void) src; (void) dst; (void) nrow; (void) n_per_row; (void) quant_weights; GGML_ABORT("TQ2P_Q4K6 quantization is done in Python (sasori)"); } +void dequantize_row_tq3p_q4k6(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { q4k6_dequantize((const uint8_t *) x, y, k, 3); } +void quantize_row_tq3p_q4k6_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { (void) x; (void) y; (void) k; GGML_ABORT("TQ3P_Q4K6 quantization is done in Python (sasori)"); } +size_t quantize_tq3p_q4k6(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { (void) src; (void) dst; (void) nrow; (void) n_per_row; (void) quant_weights; GGML_ABORT("TQ3P_Q4K6 quantization is done in Python (sasori)"); } + // ====================== "True" 2-bit (de)-quantization void dequantize_row_iq2_xxs(const block_iq2_xxs * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { @@ -5316,6 +5543,22 @@ static bool validate_e_e8m0(uint8_t e, size_t i) { } \ } +#define VALIDATE_ROW_DATA_DMN_F16_IMPL(type, data, nb, d, m, n) \ + const type * q = (const type *) (data); \ + for (size_t i = 0; i < (nb); ++i) { \ + if (!validate_fp16(q[i].d, i) || !validate_fp16(q[i].m, i) || !validate_fp16(q[i].n, i)) { \ + return false; \ + } \ + } + +#define VALIDATE_ROW_DATA_DMNO_F16_IMPL(type, data, nb, d, m, n, o) \ + const type * q = (const type *) (data); \ + for (size_t i = 0; i < (nb); ++i) { \ + if (!validate_fp16(q[i].d, i) || !validate_fp16(q[i].m, i) || !validate_fp16(q[i].n, i) || !validate_fp16(q[i].o, i)) { \ + return false; \ + } \ + } + #define VALIDATE_ROW_DATA_E_E8M0_IMPL(type, data, nb) \ const type * q = (const type *) (data); \ for (size_t i = 0; i < (nb); ++i) { \ @@ -5334,6 +5577,26 @@ static bool validate_e_e8m0(uint8_t e, size_t i) { } \ } +static bool validate_tqkp_gv(const void * data, size_t nb, int K, int g) { + const int nsub = 256 / g, qb = 64, bpb = K*qb + K*nsub*2, nscales = K*nsub; + const uint8_t * p = (const uint8_t *) data; + for (size_t i = 0; i < nb; ++i) { + const ggml_half * d = (const ggml_half *)(p + i*bpb + K*qb); + for (int j = 0; j < nscales; ++j) { if (!validate_fp16(d[j], i)) { return false; } } + } + return true; +} + +static bool validate_q4k6(const void * data, size_t nb, int K) { // g=32; only the fp16 super needs checking + const int nsub = 256 / 32, qb = 64, bpb = K*qb + K*nsub + K*2; + const uint8_t * p = (const uint8_t *) data; + for (size_t i = 0; i < nb; ++i) { + const ggml_half * d = (const ggml_half *)(p + i*bpb + K*qb + K*nsub); // K fp16 super-scales + for (int k = 0; k < K; ++k) { if (!validate_fp16(d[k], i)) { return false; } } + } + return true; +} + bool ggml_validate_row_data(enum ggml_type type, const void * data, size_t nbytes) { if (type < 0 || type >= GGML_TYPE_COUNT) { fprintf(stderr, "%s: invalid type %d\n", __func__, type); @@ -5528,6 +5791,38 @@ bool ggml_validate_row_data(enum ggml_type type, const void * data, size_t nbyte { VALIDATE_ROW_DATA_D_F16_IMPL(block_tq2_0, data, nb); } break; + case GGML_TYPE_TQ2P: + { + VALIDATE_ROW_DATA_DM_F16_IMPL(block_tq2p, data, nb, d1, d2); + } break; + case GGML_TYPE_TQ3P: + { + VALIDATE_ROW_DATA_DMN_F16_IMPL(block_tq3p, data, nb, d1, d2, d3); + } break; + case GGML_TYPE_TQ1P: + { + VALIDATE_ROW_DATA_D_F16_IMPL(block_tq1p, data, nb); + } break; + case GGML_TYPE_TQ4P: + { + VALIDATE_ROW_DATA_DMNO_F16_IMPL(block_tq4p, data, nb, d1, d2, d3, d4); + } break; + case GGML_TYPE_TQ1P_G32: { if (!validate_tqkp_gv(data, nb, 1, 32)) { return false; } } break; + case GGML_TYPE_TQ1P_G64: { if (!validate_tqkp_gv(data, nb, 1, 64)) { return false; } } break; + case GGML_TYPE_TQ1P_G128: { if (!validate_tqkp_gv(data, nb, 1, 128)) { return false; } } break; + case GGML_TYPE_TQ2P_G32: { if (!validate_tqkp_gv(data, nb, 2, 32)) { return false; } } break; + case GGML_TYPE_TQ2P_G64: { if (!validate_tqkp_gv(data, nb, 2, 64)) { return false; } } break; + case GGML_TYPE_TQ2P_G128: { if (!validate_tqkp_gv(data, nb, 2, 128)) { return false; } } break; + case GGML_TYPE_TQ3P_G32: { if (!validate_tqkp_gv(data, nb, 3, 32)) { return false; } } break; + case GGML_TYPE_TQ3P_G64: { if (!validate_tqkp_gv(data, nb, 3, 64)) { return false; } } break; + case GGML_TYPE_TQ3P_G128: { if (!validate_tqkp_gv(data, nb, 3, 128)) { return false; } } break; + case GGML_TYPE_TQ4P_G32: { if (!validate_tqkp_gv(data, nb, 4, 32)) { return false; } } break; + case GGML_TYPE_TQ4P_G64: { if (!validate_tqkp_gv(data, nb, 4, 64)) { return false; } } break; + case GGML_TYPE_TQ4P_G128: { if (!validate_tqkp_gv(data, nb, 4, 128)) { return false; } } break; + case GGML_TYPE_TQ1P_Q4K6: { if (!validate_q4k6(data, nb, 1)) { return false; } } break; + case GGML_TYPE_TQ2P_Q4K6: { if (!validate_q4k6(data, nb, 2)) { return false; } } break; + case GGML_TYPE_TQ3P_Q4K6: { if (!validate_q4k6(data, nb, 3)) { return false; } } break; + case GGML_TYPE_IQ1_S: { VALIDATE_ROW_DATA_D_F16_IMPL(block_iq1_s, data, nb); diff --git a/src/ggml-quants.h b/src/ggml-quants.h index d56c86da..ad8d927e 100644 --- a/src/ggml-quants.h +++ b/src/ggml-quants.h @@ -34,6 +34,26 @@ GGML_API void quantize_row_q8_K_ref(const float * GGML_RESTRICT x, block_q8_K * GGML_API void quantize_row_tq1_0_ref(const float * GGML_RESTRICT x, block_tq1_0 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_tq2_0_ref(const float * GGML_RESTRICT x, block_tq2_0 * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_tq2p_ref(const float * GGML_RESTRICT x, block_tq2p * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_tq3p_ref(const float * GGML_RESTRICT x, block_tq3p * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_tq1p_ref(const float * GGML_RESTRICT x, block_tq1p * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_tq1pa_ref(const float * GGML_RESTRICT x, block_tq1pa * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_tq4p_ref(const float * GGML_RESTRICT x, block_tq4p * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_tq1p_g32_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_tq1p_g64_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_tq1p_g128_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_tq2p_g32_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_tq2p_g64_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_tq2p_g128_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_tq3p_g32_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_tq3p_g64_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_tq3p_g128_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_tq4p_g32_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_tq4p_g64_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_tq4p_g128_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_tq1p_q4k6_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_tq2p_q4k6_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_tq3p_q4k6_ref(const void * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_iq3_xxs_ref(const float * GGML_RESTRICT x, block_iq3_xxs * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_iq4_nl_ref (const float * GGML_RESTRICT x, block_iq4_nl * GGML_RESTRICT y, int64_t k); @@ -62,6 +82,26 @@ GGML_API void dequantize_row_q8_K(const block_q8_K * GGML_RESTRICT x, float * GG GGML_API void dequantize_row_tq1_0(const block_tq1_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_tq2_0(const block_tq2_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq2p(const block_tq2p * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq3p(const block_tq3p * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq1p(const block_tq1p * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq1pa(const block_tq1pa * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq4p(const block_tq4p * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq1p_g32(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq1p_g64(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq1p_g128(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq2p_g32(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq2p_g64(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq2p_g128(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq3p_g32(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq3p_g64(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq3p_g128(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq4p_g32(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq4p_g64(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq4p_g128(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq1p_q4k6(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq2p_q4k6(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_tq3p_q4k6(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_iq2_xxs(const block_iq2_xxs * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_iq2_xs (const block_iq2_xs * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); @@ -86,6 +126,25 @@ GGML_API size_t quantize_iq3_s (const float * GGML_RESTRICT src, void * GGML_RE GGML_API size_t quantize_tq1_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_tq2_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_tq2p(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_tq3p(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_tq1p(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_tq4p(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_tq1p_g32(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_tq1p_g64(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_tq1p_g128(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_tq2p_g32(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_tq2p_g64(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_tq2p_g128(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_tq3p_g32(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_tq3p_g64(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_tq3p_g128(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_tq4p_g32(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_tq4p_g64(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_tq4p_g128(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_tq1p_q4k6(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_tq2p_q4k6(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_tq3p_q4k6(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_K(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_q3_K(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); diff --git a/src/ggml.c b/src/ggml.c index 0f682fd1..ae5e3c72 100644 --- a/src/ggml.c +++ b/src/ggml.c @@ -906,6 +906,166 @@ static const struct ggml_type_traits type_traits[GGML_TYPE_COUNT] = { .to_float = (ggml_to_float_t) dequantize_row_tq2_0, .from_float_ref = (ggml_from_float_t) quantize_row_tq2_0_ref, }, + [GGML_TYPE_TQ2P] = { + .type_name = "tq2p", + .blck_size = QK_TQ2P, + .type_size = sizeof(block_tq2p), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq2p, + .from_float_ref = (ggml_from_float_t) quantize_row_tq2p_ref, + }, + [GGML_TYPE_TQ3P] = { + .type_name = "tq3p", + .blck_size = QK_TQ3P, + .type_size = sizeof(block_tq3p), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq3p, + .from_float_ref = (ggml_from_float_t) quantize_row_tq3p_ref, + }, + [GGML_TYPE_TQ1P] = { + .type_name = "tq1p", + .blck_size = QK_TQ1P, + .type_size = sizeof(block_tq1p), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq1p, + .from_float_ref = (ggml_from_float_t) quantize_row_tq1p_ref, + }, + [GGML_TYPE_TQ4P] = { + .type_name = "tq4p", + .blck_size = QK_TQ4P, + .type_size = sizeof(block_tq4p), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq4p, + .from_float_ref = (ggml_from_float_t) quantize_row_tq4p_ref, + }, + [GGML_TYPE_TQ1P_G32] = { + .type_name = "tq1p_g32", + .blck_size = QK_TQKP_SUPER, + .type_size = TQKP_GV_BYTES(1, 32), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq1p_g32, + .from_float_ref = (ggml_from_float_t) quantize_row_tq1p_g32_ref, + }, + [GGML_TYPE_TQ1P_G64] = { + .type_name = "tq1p_g64", + .blck_size = QK_TQKP_SUPER, + .type_size = TQKP_GV_BYTES(1, 64), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq1p_g64, + .from_float_ref = (ggml_from_float_t) quantize_row_tq1p_g64_ref, + }, + [GGML_TYPE_TQ1P_G128] = { + .type_name = "tq1p_g128", + .blck_size = QK_TQKP_SUPER, + .type_size = TQKP_GV_BYTES(1, 128), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq1p_g128, + .from_float_ref = (ggml_from_float_t) quantize_row_tq1p_g128_ref, + }, + [GGML_TYPE_TQ2P_G32] = { + .type_name = "tq2p_g32", + .blck_size = QK_TQKP_SUPER, + .type_size = TQKP_GV_BYTES(2, 32), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq2p_g32, + .from_float_ref = (ggml_from_float_t) quantize_row_tq2p_g32_ref, + }, + [GGML_TYPE_TQ2P_G64] = { + .type_name = "tq2p_g64", + .blck_size = QK_TQKP_SUPER, + .type_size = TQKP_GV_BYTES(2, 64), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq2p_g64, + .from_float_ref = (ggml_from_float_t) quantize_row_tq2p_g64_ref, + }, + [GGML_TYPE_TQ2P_G128] = { + .type_name = "tq2p_g128", + .blck_size = QK_TQKP_SUPER, + .type_size = TQKP_GV_BYTES(2, 128), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq2p_g128, + .from_float_ref = (ggml_from_float_t) quantize_row_tq2p_g128_ref, + }, + [GGML_TYPE_TQ3P_G32] = { + .type_name = "tq3p_g32", + .blck_size = QK_TQKP_SUPER, + .type_size = TQKP_GV_BYTES(3, 32), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq3p_g32, + .from_float_ref = (ggml_from_float_t) quantize_row_tq3p_g32_ref, + }, + [GGML_TYPE_TQ3P_G64] = { + .type_name = "tq3p_g64", + .blck_size = QK_TQKP_SUPER, + .type_size = TQKP_GV_BYTES(3, 64), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq3p_g64, + .from_float_ref = (ggml_from_float_t) quantize_row_tq3p_g64_ref, + }, + [GGML_TYPE_TQ3P_G128] = { + .type_name = "tq3p_g128", + .blck_size = QK_TQKP_SUPER, + .type_size = TQKP_GV_BYTES(3, 128), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq3p_g128, + .from_float_ref = (ggml_from_float_t) quantize_row_tq3p_g128_ref, + }, + [GGML_TYPE_TQ4P_G32] = { + .type_name = "tq4p_g32", + .blck_size = QK_TQKP_SUPER, + .type_size = TQKP_GV_BYTES(4, 32), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq4p_g32, + .from_float_ref = (ggml_from_float_t) quantize_row_tq4p_g32_ref, + }, + [GGML_TYPE_TQ4P_G64] = { + .type_name = "tq4p_g64", + .blck_size = QK_TQKP_SUPER, + .type_size = TQKP_GV_BYTES(4, 64), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq4p_g64, + .from_float_ref = (ggml_from_float_t) quantize_row_tq4p_g64_ref, + }, + [GGML_TYPE_TQ4P_G128] = { + .type_name = "tq4p_g128", + .blck_size = QK_TQKP_SUPER, + .type_size = TQKP_GV_BYTES(4, 128), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq4p_g128, + .from_float_ref = (ggml_from_float_t) quantize_row_tq4p_g128_ref, + }, + [GGML_TYPE_TQ1P_Q4K6] = { + .type_name = "tq1p_q4k6", + .blck_size = QK_TQKP_SUPER, + .type_size = TQKP_Q4K6_BYTES(1), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq1p_q4k6, + .from_float_ref = (ggml_from_float_t) quantize_row_tq1p_q4k6_ref, + }, + [GGML_TYPE_TQ2P_Q4K6] = { + .type_name = "tq2p_q4k6", + .blck_size = QK_TQKP_SUPER, + .type_size = TQKP_Q4K6_BYTES(2), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq2p_q4k6, + .from_float_ref = (ggml_from_float_t) quantize_row_tq2p_q4k6_ref, + }, + [GGML_TYPE_TQ3P_Q4K6] = { + .type_name = "tq3p_q4k6", + .blck_size = QK_TQKP_SUPER, + .type_size = TQKP_Q4K6_BYTES(3), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq3p_q4k6, + .from_float_ref = (ggml_from_float_t) quantize_row_tq3p_q4k6_ref, + }, + [GGML_TYPE_TQ1PA] = { + .type_name = "tq1pa", + .blck_size = QK_TQ1P, + .type_size = sizeof(block_tq1pa), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_tq1pa, + .from_float_ref = (ggml_from_float_t) quantize_row_tq1pa_ref, + }, [36] = { // GGML_TYPE_IQ4_NL_4_4 .type_name = "TYPE_IQ4_NL_4_4 REMOVED, use IQ4_NL with runtime repacking", .blck_size = 0, @@ -7743,6 +7903,25 @@ size_t ggml_quantize_chunk( case GGML_TYPE_Q6_K: result = quantize_q6_K (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_TQ1_0: result = quantize_tq1_0 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_TQ2_0: result = quantize_tq2_0 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_TQ2P: result = quantize_tq2p (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_TQ3P: result = quantize_tq3p (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_TQ1P: result = quantize_tq1p (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_TQ4P: result = quantize_tq4p (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_TQ1P_G32: result = quantize_tq1p_g32(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_TQ1P_G64: result = quantize_tq1p_g64(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_TQ1P_G128: result = quantize_tq1p_g128(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_TQ2P_G32: result = quantize_tq2p_g32(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_TQ2P_G64: result = quantize_tq2p_g64(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_TQ2P_G128: result = quantize_tq2p_g128(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_TQ3P_G32: result = quantize_tq3p_g32(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_TQ3P_G64: result = quantize_tq3p_g64(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_TQ3P_G128: result = quantize_tq3p_g128(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_TQ4P_G32: result = quantize_tq4p_g32(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_TQ4P_G64: result = quantize_tq4p_g64(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_TQ4P_G128: result = quantize_tq4p_g128(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_TQ1P_Q4K6: result = quantize_tq1p_q4k6(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_TQ2P_Q4K6: result = quantize_tq2p_q4k6(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_TQ3P_Q4K6: result = quantize_tq3p_q4k6(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_IQ2_XXS: result = quantize_iq2_xxs(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_IQ2_XS: result = quantize_iq2_xs (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_IQ3_XXS: result = quantize_iq3_xxs(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;