diff --git "a/kernel/sdcpp-cpu.patch" "b/kernel/sdcpp-cpu.patch" new file mode 100644--- /dev/null +++ "b/kernel/sdcpp-cpu.patch" @@ -0,0 +1,1783 @@ +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;