summaryrefslogtreecommitdiff
path: root/ggml-quants.c
diff options
context:
space:
mode:
authorKawrakow <48489457+ikawrakow@users.noreply.github.com>2024-02-24 16:23:52 +0200
committerGitHub <noreply@github.com>2024-02-24 16:23:52 +0200
commit4c4cb30736582cacb1a164a9d4bc8e17b1014be7 (patch)
tree5f953370b3124531d9cbb9b9d5cfdb264ddf60bc /ggml-quants.c
parent525213d2f5da1eaf4b922b6b792cb52b2c613368 (diff)
IQ3_S: a much better alternative to Q3_K (#5676)
* iq4_nl: squash commits for easier rebase * Basics (quantize, dequantize) * CUDA dequantize and dot product * Slightly faster CUDA dot product (120 t/s) * Switch to 6-bit scales * Scalar dot product * AVX2 dot product * ARM_NEON dot product * Works on metal, but still slow * Slightly better Metal dot product * Another small Metal improvement * Metal dot product is getting there * Faster CUDA dot product * Add 1/8 ffn_down layers as Q5_K when no imatrix has been provided * Report the actual bpw * Add _xs mix that is 4.05 bpw for non-MoE models * Remove IQ4_XS for now, slightly adjust kvalues_iq4nl * AVX2 dot product uses Q8_0 instead of Q8_K * Add to test-backend-ops * Minor fix * Also use use Q5_K for attn_output in MoE models * Fixes after merging latest master * Switching to blocks of 32 * AVX2 for blocks of 32 * Scaler dot product for blocks of 32 * ARM_NEON dot product for blocks of 32 * Metal kernels for blocks of 32 * Slightly faster Metal kernels * Resurrecting iq3_xs After all the experimentation, nothing was better than this. * Minor PPL improvement via a block scale fudge factor * Minor improvement via 3 neighbours * iq3_xs: working scalar and AVX2 dot products * iq3_xs: ARM_NEON dot product - works but extremely slow (10 t/s) * iq3_xs: working Metal implementation * Adding IQ3_M - IQ3_XS mix with mostly Q4_K * iiq3_xs: a 3.4375 bpw variant * iq3_xs: make CUDA work for new version * iq3_xs: make scalar and AVX2 work for new version * iq3_s: make ARM_NEON work with new version * iq3_xs: make new version work on metal Performance is very similar to Q3_K_S * iq3_xs: tiny Metal speed improvement * iq3_xs: tiny Metal speed improvement * Fix stupid warning * Q3_K_XS now uses a mix of IQ3_XS and IQ3_XXS * iq3_xs: rename to iq3_s * iq3_s: make tests pass * Move Q3_K_XS mix to 3.25 bpw * Attempt to fix failing tests * Another attempt to fix the Windows builds * Attempt to fix ROCm * ROCm again * iq3_s: partial fix for QK_K = 64 * iq3_s: make it work on metal for QK_K = 64 Pleasent surprise: the coding was super-block size independent, so all it took was to delete some QK_K == 256 guards. * Will this fix ROCm? --------- Co-authored-by: Iwan Kawrakow <iwan.kawrakow@gmail.com>
Diffstat (limited to 'ggml-quants.c')
-rw-r--r--ggml-quants.c674
1 files changed, 610 insertions, 64 deletions
diff --git a/ggml-quants.c b/ggml-quants.c
index b15977f5..5c5f2ce1 100644
--- a/ggml-quants.c
+++ b/ggml-quants.c
@@ -3505,6 +3505,73 @@ static const uint32_t iq3xxs_grid[256] = {
0x3e1c1c1c, 0x3e1c3404, 0x3e24140c, 0x3e24240c, 0x3e2c0404, 0x3e2c0414, 0x3e2c1424, 0x3e341c04,
};
+static const uint32_t iq3xs_grid[512] = {
+ 0x04040404, 0x0404040c, 0x04040414, 0x0404042c, 0x0404043e, 0x04040c04, 0x04040c0c, 0x04040c14,
+ 0x04040c24, 0x04040c34, 0x04041404, 0x0404140c, 0x0404142c, 0x04041c1c, 0x04042404, 0x04042414,
+ 0x0404242c, 0x0404243e, 0x04042c0c, 0x04042c1c, 0x04043404, 0x04043414, 0x04043e0c, 0x04043e24,
+ 0x04043e3e, 0x040c0404, 0x040c040c, 0x040c0414, 0x040c0424, 0x040c0c04, 0x040c0c0c, 0x040c0c2c,
+ 0x040c1404, 0x040c141c, 0x040c143e, 0x040c1c0c, 0x040c1c2c, 0x040c2424, 0x040c340c, 0x040c342c,
+ 0x040c3e14, 0x04140404, 0x0414040c, 0x0414042c, 0x0414043e, 0x04140c04, 0x04140c1c, 0x04140c34,
+ 0x0414140c, 0x0414142c, 0x04141c04, 0x04141c24, 0x04142414, 0x0414242c, 0x0414243e, 0x04142c0c,
+ 0x04142c1c, 0x04143e04, 0x04143e1c, 0x041c041c, 0x041c0c0c, 0x041c0c2c, 0x041c1404, 0x041c1414,
+ 0x041c1c0c, 0x041c1c1c, 0x041c1c34, 0x041c2424, 0x041c2c04, 0x041c2c14, 0x041c343e, 0x041c3e0c,
+ 0x041c3e2c, 0x04240404, 0x04240c1c, 0x04240c3e, 0x0424140c, 0x04241424, 0x04241c14, 0x04242404,
+ 0x0424241c, 0x04242c0c, 0x04243e04, 0x042c0414, 0x042c0424, 0x042c1404, 0x042c1414, 0x042c1434,
+ 0x042c1c1c, 0x042c240c, 0x042c242c, 0x042c243e, 0x042c3434, 0x042c3e1c, 0x04340434, 0x04340c0c,
+ 0x04340c1c, 0x04341c0c, 0x04342c14, 0x04343e0c, 0x043e0404, 0x043e0414, 0x043e0424, 0x043e1404,
+ 0x043e1414, 0x043e1434, 0x043e1c1c, 0x043e2c04, 0x043e2c24, 0x0c040404, 0x0c04040c, 0x0c040414,
+ 0x0c040424, 0x0c040c04, 0x0c040c0c, 0x0c040c1c, 0x0c040c2c, 0x0c040c3e, 0x0c041404, 0x0c041414,
+ 0x0c041c0c, 0x0c041c24, 0x0c041c34, 0x0c042c24, 0x0c042c34, 0x0c04340c, 0x0c043e14, 0x0c0c0404,
+ 0x0c0c040c, 0x0c0c041c, 0x0c0c0434, 0x0c0c0c04, 0x0c0c0c24, 0x0c0c140c, 0x0c0c1c04, 0x0c0c1c1c,
+ 0x0c0c240c, 0x0c0c2c04, 0x0c0c2c14, 0x0c0c3e04, 0x0c0c3e34, 0x0c140404, 0x0c140c14, 0x0c140c2c,
+ 0x0c140c3e, 0x0c141404, 0x0c141424, 0x0c141c14, 0x0c142404, 0x0c14241c, 0x0c142c2c, 0x0c143404,
+ 0x0c143e14, 0x0c1c040c, 0x0c1c0424, 0x0c1c043e, 0x0c1c0c04, 0x0c1c0c1c, 0x0c1c140c, 0x0c1c143e,
+ 0x0c1c1c04, 0x0c1c1c24, 0x0c1c240c, 0x0c1c3414, 0x0c1c3e04, 0x0c24041c, 0x0c24042c, 0x0c240c14,
+ 0x0c240c24, 0x0c241c0c, 0x0c241c1c, 0x0c242414, 0x0c242434, 0x0c242c04, 0x0c242c24, 0x0c2c040c,
+ 0x0c2c0c04, 0x0c2c0c1c, 0x0c2c140c, 0x0c2c1c04, 0x0c2c1c14, 0x0c2c2c0c, 0x0c341404, 0x0c341424,
+ 0x0c34143e, 0x0c342424, 0x0c342434, 0x0c3e040c, 0x0c3e041c, 0x0c3e0c04, 0x0c3e0c14, 0x0c3e140c,
+ 0x0c3e1c2c, 0x0c3e240c, 0x0c3e3414, 0x0c3e3e04, 0x14040404, 0x1404040c, 0x1404041c, 0x1404042c,
+ 0x1404043e, 0x14040c04, 0x14040c14, 0x14040c24, 0x14040c34, 0x1404140c, 0x1404141c, 0x1404143e,
+ 0x14041c04, 0x14041c14, 0x1404240c, 0x1404241c, 0x1404242c, 0x14042c04, 0x14042c14, 0x1404343e,
+ 0x14043e04, 0x14043e1c, 0x14043e2c, 0x140c0404, 0x140c0414, 0x140c0c04, 0x140c0c1c, 0x140c0c3e,
+ 0x140c1414, 0x140c142c, 0x140c1c0c, 0x140c1c24, 0x140c2414, 0x140c2c0c, 0x1414040c, 0x14140424,
+ 0x1414043e, 0x1414140c, 0x1414141c, 0x14141c04, 0x14141c3e, 0x1414240c, 0x14142c1c, 0x14142c3e,
+ 0x14143e0c, 0x14143e24, 0x141c0404, 0x141c0414, 0x141c042c, 0x141c0c0c, 0x141c1414, 0x141c1424,
+ 0x141c1c0c, 0x141c1c1c, 0x141c2414, 0x141c2c04, 0x141c3434, 0x1424040c, 0x1424043e, 0x14241404,
+ 0x1424141c, 0x14241c14, 0x14241c2c, 0x1424240c, 0x14243e14, 0x14243e2c, 0x142c0424, 0x142c0c0c,
+ 0x142c1414, 0x142c1c3e, 0x142c2404, 0x142c2c1c, 0x142c3e04, 0x14340404, 0x14340414, 0x1434043e,
+ 0x1434140c, 0x14342c2c, 0x1434340c, 0x143e042c, 0x143e0c0c, 0x143e1434, 0x143e1c04, 0x143e241c,
+ 0x143e2c04, 0x1c040414, 0x1c040c0c, 0x1c040c1c, 0x1c040c2c, 0x1c040c3e, 0x1c041414, 0x1c041c0c,
+ 0x1c041c1c, 0x1c041c2c, 0x1c042414, 0x1c042424, 0x1c04243e, 0x1c042c0c, 0x1c04341c, 0x1c043e0c,
+ 0x1c0c040c, 0x1c0c041c, 0x1c0c042c, 0x1c0c0c24, 0x1c0c140c, 0x1c0c141c, 0x1c0c2404, 0x1c0c3404,
+ 0x1c0c3e14, 0x1c0c3e34, 0x1c140404, 0x1c140c14, 0x1c141404, 0x1c141c14, 0x1c141c24, 0x1c142c04,
+ 0x1c1c040c, 0x1c1c0c04, 0x1c1c0c24, 0x1c1c140c, 0x1c1c141c, 0x1c1c143e, 0x1c1c1c04, 0x1c1c240c,
+ 0x1c1c241c, 0x1c1c243e, 0x1c1c2c2c, 0x1c1c3e1c, 0x1c24041c, 0x1c240c0c, 0x1c240c34, 0x1c241414,
+ 0x1c241c0c, 0x1c242c14, 0x1c243404, 0x1c243424, 0x1c2c040c, 0x1c2c0c04, 0x1c2c0c14, 0x1c2c142c,
+ 0x1c2c1c14, 0x1c2c2424, 0x1c2c2c34, 0x1c2c3e1c, 0x1c340c34, 0x1c34240c, 0x1c3e040c, 0x1c3e041c,
+ 0x1c3e1404, 0x1c3e1414, 0x1c3e1c2c, 0x24040404, 0x24040424, 0x24040c14, 0x24041404, 0x24041424,
+ 0x2404143e, 0x24041c14, 0x2404240c, 0x24042c04, 0x24043e04, 0x240c0414, 0x240c043e, 0x240c0c0c,
+ 0x240c0c1c, 0x240c1414, 0x240c1c04, 0x240c1c2c, 0x240c241c, 0x240c2c0c, 0x240c2c2c, 0x2414040c,
+ 0x2414041c, 0x24140c04, 0x24140c2c, 0x2414140c, 0x24141c1c, 0x24142404, 0x24142c3e, 0x24143414,
+ 0x24143e04, 0x241c0424, 0x241c0c0c, 0x241c0c1c, 0x241c1404, 0x241c1414, 0x241c1c0c, 0x241c1c2c,
+ 0x24240404, 0x24240414, 0x24241424, 0x24241c3e, 0x24242404, 0x24243e0c, 0x242c042c, 0x242c043e,
+ 0x242c140c, 0x242c3414, 0x24340c1c, 0x24341c24, 0x24343404, 0x243e0c04, 0x243e0c2c, 0x243e1c04,
+ 0x243e241c, 0x243e2c0c, 0x2c040414, 0x2c040c04, 0x2c040c24, 0x2c041414, 0x2c042404, 0x2c042424,
+ 0x2c04243e, 0x2c042c14, 0x2c043434, 0x2c043e24, 0x2c0c040c, 0x2c0c041c, 0x2c0c042c, 0x2c0c0c14,
+ 0x2c0c140c, 0x2c0c1c14, 0x2c0c3e14, 0x2c140404, 0x2c140c0c, 0x2c14141c, 0x2c141c04, 0x2c141c34,
+ 0x2c142c1c, 0x2c1c0414, 0x2c1c043e, 0x2c1c0c04, 0x2c1c143e, 0x2c1c2424, 0x2c1c2c0c, 0x2c1c342c,
+ 0x2c1c3e1c, 0x2c24040c, 0x2c240424, 0x2c241404, 0x2c241c14, 0x2c242434, 0x2c2c0c14, 0x2c2c1434,
+ 0x2c2c2c0c, 0x2c2c2c1c, 0x2c342414, 0x2c3e0414, 0x2c3e0424, 0x2c3e1414, 0x34040c0c, 0x34040c1c,
+ 0x34040c2c, 0x34041c0c, 0x34041c1c, 0x34043404, 0x340c0404, 0x340c1404, 0x340c143e, 0x340c3424,
+ 0x34140c14, 0x34141c24, 0x34142414, 0x34142c2c, 0x34143414, 0x34143e04, 0x341c0404, 0x341c0c24,
+ 0x341c140c, 0x341c2404, 0x3424142c, 0x3424241c, 0x34243414, 0x342c0404, 0x342c041c, 0x342c1c24,
+ 0x342c3404, 0x3434042c, 0x34342404, 0x343e0c0c, 0x343e0c1c, 0x3e040404, 0x3e040424, 0x3e04043e,
+ 0x3e041404, 0x3e041414, 0x3e041c34, 0x3e042404, 0x3e042c24, 0x3e043414, 0x3e0c0414, 0x3e0c0c0c,
+ 0x3e0c1424, 0x3e0c241c, 0x3e0c242c, 0x3e14040c, 0x3e140424, 0x3e140c04, 0x3e140c34, 0x3e14140c,
+ 0x3e141c04, 0x3e142c0c, 0x3e1c0414, 0x3e1c1c14, 0x3e1c1c2c, 0x3e1c2c1c, 0x3e24040c, 0x3e24042c,
+ 0x3e240c1c, 0x3e241404, 0x3e242c04, 0x3e2c1414, 0x3e2c2414, 0x3e340414, 0x3e341c0c, 0x3e3e0404,
+};
+
#define NGRID_IQ2XXS 512
static const uint64_t iq1s_grid[NGRID_IQ2XXS] = {
0xffffffffffff0101, 0xffffffffff01ff00, 0xffffffffff010100, 0xffffffff00000000,
@@ -3736,6 +3803,49 @@ void dequantize_row_iq3_xxs(const block_iq3_xxs * restrict x, float * restrict y
}
}
+// ====================== 3.3125 bpw (de)-quantization
+
+void dequantize_row_iq3_s(const block_iq3_s * restrict x, float * restrict y, int k) {
+ assert(k % QK_K == 0);
+ const int nb = k / QK_K;
+
+ for (int i = 0; i < nb; i++) {
+
+ const float d = GGML_FP16_TO_FP32(x[i].d);
+ const uint8_t * qs = x[i].qs;
+ const uint8_t * qh = x[i].qh;
+ const uint8_t * signs = x[i].signs;
+
+ for (int ib32 = 0; ib32 < QK_K/32; ib32 += 2) {
+ const float db1 = d * (0.5f + (x[i].scales[ib32/2] & 0xf)) * 0.5f;
+ const float db2 = d * (0.5f + (x[i].scales[ib32/2] >> 4)) * 0.5f;
+ for (int l = 0; l < 4; ++l) {
+ const uint8_t * grid1 = (const uint8_t *)(iq3xs_grid + (qs[2*l+0] | ((qh[0] << (8-2*l)) & 256)));
+ const uint8_t * grid2 = (const uint8_t *)(iq3xs_grid + (qs[2*l+1] | ((qh[0] << (7-2*l)) & 256)));
+ for (int j = 0; j < 4; ++j) {
+ y[j+0] = db1 * grid1[j] * (signs[l] & kmask_iq2xs[j+0] ? -1.f : 1.f);
+ y[j+4] = db1 * grid2[j] * (signs[l] & kmask_iq2xs[j+4] ? -1.f : 1.f);
+ }
+ y += 8;
+ }
+ qs += 8;
+ signs += 4;
+ for (int l = 0; l < 4; ++l) {
+ const uint8_t * grid1 = (const uint8_t *)(iq3xs_grid + (qs[2*l+0] | ((qh[1] << (8-2*l)) & 256)));
+ const uint8_t * grid2 = (const uint8_t *)(iq3xs_grid + (qs[2*l+1] | ((qh[1] << (7-2*l)) & 256)));
+ for (int j = 0; j < 4; ++j) {
+ y[j+0] = db2 * grid1[j] * (signs[l] & kmask_iq2xs[j+0] ? -1.f : 1.f);
+ y[j+4] = db2 * grid2[j] * (signs[l] & kmask_iq2xs[j+4] ? -1.f : 1.f);
+ }
+ y += 8;
+ }
+ qh += 2;
+ qs += 8;
+ signs += 4;
+ }
+ }
+}
+
// ====================== 1.5625 bpw (de)-quantization
void dequantize_row_iq1_s(const block_iq1_s * restrict x, float * restrict y, int k) {
@@ -8806,6 +8916,7 @@ void ggml_vec_dot_q6_K_q8_K(int n, float * restrict s, size_t bs, const void * r
#endif
+#if defined (__AVX2__) || defined (__ARM_NEON)
static const int8_t keven_signs_q2xs[1024] = {
1, 1, 1, 1, 1, 1, 1, 1, -1, 1, 1, 1, 1, 1, 1, -1, 1, -1, 1, 1, 1, 1, 1, -1, -1, -1, 1, 1, 1, 1, 1, 1,
1, 1, -1, 1, 1, 1, 1, -1, -1, 1, -1, 1, 1, 1, 1, 1, 1, -1, -1, 1, 1, 1, 1, 1, -1, -1, -1, 1, 1, 1, 1, -1,
@@ -8840,6 +8951,7 @@ static const int8_t keven_signs_q2xs[1024] = {
1, 1, 1, -1, -1, -1, -1, 1, -1, 1, 1, -1, -1, -1, -1, -1, 1, -1, 1, -1, -1, -1, -1, -1, -1, -1, 1, -1, -1, -1, -1, 1,
1, 1, -1, -1, -1, -1, -1, -1, -1, 1, -1, -1, -1, -1, -1, 1, 1, -1, -1, -1, -1, -1, -1, 1, -1, -1, -1, -1, -1, -1, -1, -1,
};
+#endif
void ggml_vec_dot_iq2_xxs_q8_K(int n, float * restrict s, size_t bs, const void * restrict vx, size_t bx, const void * restrict vy, size_t by, int nrc) {
assert(n % QK_K == 0);
@@ -9327,6 +9439,202 @@ void ggml_vec_dot_iq3_xxs_q8_K(int n, float * restrict s, size_t bs, const void
#endif
}
+void ggml_vec_dot_iq3_s_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_K == 0);
+ assert(nrc == 1);
+ UNUSED(nrc);
+ UNUSED(bx);
+ UNUSED(by);
+ UNUSED(bs);
+
+ const block_iq3_s * restrict x = vx;
+ const block_q8_K * restrict y = vy;
+
+ const int nb = n / QK_K;
+
+#if defined(__ARM_NEON)
+
+ static const uint8_t k_mask1[32] = {0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01,
+ 0x02, 0x02, 0x02, 0x02, 0x02, 0x02, 0x02, 0x02, 0x03, 0x03, 0x03, 0x03, 0x03, 0x03, 0x03, 0x03
+ };
+
+ static const uint8_t k_mask2[16] = {0x01, 0x02, 0x04, 0x08, 0x10, 0x20, 0x40, 0x80, 0x01, 0x02, 0x04, 0x08, 0x10, 0x20, 0x40, 0x80,};
+
+ const uint8x16x2_t mask1 = vld1q_u8_x2(k_mask1);
+ const uint8x16_t mask2 = vld1q_u8(k_mask2);
+
+ uint8x16x2_t vs;
+ ggml_int8x16x4_t q3s;
+ ggml_int8x16x4_t q8b;
+
+ float sumf = 0;
+ for (int i = 0; i < nb; ++i) {
+ const float d = GGML_FP16_TO_FP32(x[i].d) * y[i].d;
+ const uint8_t * restrict qs = x[i].qs;
+ const uint8_t * restrict qh = x[i].qh;
+ const uint16_t * restrict signs = (const uint16_t *)x[i].signs;
+ const int8_t * restrict q8 = y[i].qs;
+ int sumi1 = 0, sumi2 = 0;
+ for (int ib32 = 0; ib32 < QK_K/32; ib32 += 2) {
+ q8b = ggml_vld1q_s8_x4(q8); q8 += 64;
+ const uint32x4_t aux32x4_0 = {iq3xs_grid[qs[ 0] | ((qh[ib32+0] << 8) & 256)], iq3xs_grid[qs[ 1] | ((qh[ib32+0] << 7) & 256)],
+ iq3xs_grid[qs[ 2] | ((qh[ib32+0] << 6) & 256)], iq3xs_grid[qs[ 3] | ((qh[ib32+0] << 5) & 256)]};
+ const uint32x4_t aux32x4_1 = {iq3xs_grid[qs[ 4] | ((qh[ib32+0] << 4) & 256)], iq3xs_grid[qs[ 5] | ((qh[ib32+0] << 3) & 256)],
+ iq3xs_grid[qs[ 6] | ((qh[ib32+0] << 2) & 256)], iq3xs_grid[qs[ 7] | ((qh[ib32+0] << 1) & 256)]};
+ const uint32x4_t aux32x4_2 = {iq3xs_grid[qs[ 8] | ((qh[ib32+1] << 8) & 256)], iq3xs_grid[qs[ 9] | ((qh[ib32+1] << 7) & 256)],
+ iq3xs_grid[qs[10] | ((qh[ib32+1] << 6) & 256)], iq3xs_grid[qs[11] | ((qh[ib32+1] << 5) & 256)]};
+ const uint32x4_t aux32x4_3 = {iq3xs_grid[qs[12] | ((qh[ib32+1] << 4) & 256)], iq3xs_grid[qs[13] | ((qh[ib32+1] << 3) & 256)],
+ iq3xs_grid[qs[14] | ((qh[ib32+1] << 2) & 256)], iq3xs_grid[qs[15] | ((qh[ib32+1] << 1) & 256)]};
+ qs += 16;
+
+ vs.val[0] = vreinterpretq_u8_u32(vdupq_n_u32(signs[0] | (signs[1] << 16)));
+ vs.val[1] = vandq_u8(vqtbl1q_u8(vs.val[0], mask1.val[1]), mask2);
+ vs.val[0] = vandq_u8(vqtbl1q_u8(vs.val[0], mask1.val[0]), mask2);
+ vs.val[0] = vceqq_u8(vs.val[0], mask2);
+ vs.val[1] = vceqq_u8(vs.val[1], mask2);
+
+ q3s.val[0] = vsubq_s8(vreinterpretq_s8_u8(veorq_u8(vs.val[0], vreinterpretq_u8_u32(aux32x4_0))), vreinterpretq_s8_u8(vs.val[0]));
+ q3s.val[1] = vsubq_s8(vreinterpretq_s8_u8(veorq_u8(vs.val[1], vreinterpretq_u8_u32(aux32x4_1))), vreinterpretq_s8_u8(vs.val[1]));
+
+ vs.val[0] = vreinterpretq_u8_u32(vdupq_n_u32(signs[2] | (signs[3] << 16)));
+ vs.val[1] = vandq_u8(vqtbl1q_u8(vs.val[0], mask1.val[1]), mask2);
+ vs.val[0] = vandq_u8(vqtbl1q_u8(vs.val[0], mask1.val[0]), mask2);
+ vs.val[0] = vceqq_u8(vs.val[0], mask2);
+ vs.val[1] = vceqq_u8(vs.val[1], mask2);
+
+ signs += 4;
+
+ q3s.val[2] = vsubq_s8(vreinterpretq_s8_u8(veorq_u8(vs.val[0], vreinterpretq_u8_u32(aux32x4_2))), vreinterpretq_s8_u8(vs.val[0]));
+ q3s.val[3] = vsubq_s8(vreinterpretq_s8_u8(veorq_u8(vs.val[1], vreinterpretq_u8_u32(aux32x4_3))), vreinterpretq_s8_u8(vs.val[1]));
+
+ const int32x4_t p1 = ggml_vdotq_s32(ggml_vdotq_s32(vdupq_n_s32(0), q3s.val[0], q8b.val[0]), q3s.val[1], q8b.val[1]);
+ const int32x4_t p2 = ggml_vdotq_s32(ggml_vdotq_s32(vdupq_n_s32(0), q3s.val[2], q8b.val[2]), q3s.val[3], q8b.val[3]);
+ sumi1 += vaddvq_s32(p1) * (1 + 2*(x[i].scales[ib32/2] & 0xf));
+ sumi2 += vaddvq_s32(p2) * (1 + 2*(x[i].scales[ib32/2] >> 4));
+ }
+ sumf += d*(sumi1 + sumi2);
+ }
+ *s = 0.25f * sumf;
+
+#elif defined(__AVX2__)
+
+ static const uint8_t k_mask1[32] = {0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01,
+ 0x02, 0x02, 0x02, 0x02, 0x02, 0x02, 0x02, 0x02, 0x03, 0x03, 0x03, 0x03, 0x03, 0x03, 0x03, 0x03
+ };
+
+ static const uint8_t k_mask2[32] = {0x01, 0x02, 0x04, 0x08, 0x10, 0x20, 0x40, 0x80, 0x01, 0x02, 0x04, 0x08, 0x10, 0x20, 0x40, 0x80,
+ 0x01, 0x02, 0x04, 0x08, 0x10, 0x20, 0x40, 0x80, 0x01, 0x02, 0x04, 0x08, 0x10, 0x20, 0x40, 0x80,
+ };
+
+ const __m256i mask1 = _mm256_loadu_si256((const __m256i*)k_mask1);
+ const __m256i mask2 = _mm256_loadu_si256((const __m256i*)k_mask2);
+
+ __m256 accumf = _mm256_setzero_ps();
+ for (int i = 0; i < nb; ++i) {
+ const float d = GGML_FP16_TO_FP32(x[i].d) * y[i].d;
+ const uint8_t * restrict qs = x[i].qs;
+ const uint8_t * restrict qh = x[i].qh;
+ const uint16_t * restrict signs = (const uint16_t *)x[i].signs;
+ const int8_t * restrict q8 = y[i].qs;
+ __m256i sumi1 = _mm256_setzero_si256();
+ __m256i sumi2 = _mm256_setzero_si256();
+ for (int ib32 = 0; ib32 < QK_K/32; ib32 += 2) {
+ const __m256i q8_1 = _mm256_loadu_si256((const __m256i *)q8); q8 += 32;
+ const __m256i q8_2 = _mm256_loadu_si256((const __m256i *)q8); q8 += 32;
+ const __m256i q2_1 = _mm256_set_epi32(iq3xs_grid[qs[7] | ((qh[ib32+0] << 1) & 256)],
+ iq3xs_grid[qs[6] | ((qh[ib32+0] << 2) & 256)],
+ iq3xs_grid[qs[5] | ((qh[ib32+0] << 3) & 256)],
+ iq3xs_grid[qs[4] | ((qh[ib32+0] << 4) & 256)],
+ iq3xs_grid[qs[3] | ((qh[ib32+0] << 5) & 256)],
+ iq3xs_grid[qs[2] | ((qh[ib32+0] << 6) & 256)],
+ iq3xs_grid[qs[1] | ((qh[ib32+0] << 7) & 256)],
+ iq3xs_grid[qs[0] | ((qh[ib32+0] << 8) & 256)]);
+ qs += 8;
+ const __m256i q2_2 = _mm256_set_epi32(iq3xs_grid[qs[7] | ((qh[ib32+1] << 1) & 256)],
+ iq3xs_grid[qs[6] | ((qh[ib32+1] << 2) & 256)],
+ iq3xs_grid[qs[5] | ((qh[ib32+1] << 3) & 256)],
+ iq3xs_grid[qs[4] | ((qh[ib32+1] << 4) & 256)],
+ iq3xs_grid[qs[3] | ((qh[ib32+1] << 5) & 256)],
+ iq3xs_grid[qs[2] | ((qh[ib32+1] << 6) & 256)],
+ iq3xs_grid[qs[1] | ((qh[ib32+1] << 7) & 256)],
+ iq3xs_grid[qs[0] | ((qh[ib32+1] << 8) & 256)]);
+ qs += 8;
+
+ __m256i aux256 = _mm256_set1_epi32(signs[0] | (signs[1] << 16));
+ aux256 = _mm256_and_si256(_mm256_shuffle_epi8(aux256,mask1), mask2);
+ const __m256i s2_1 = _mm256_cmpeq_epi8(aux256, mask2);
+ const __m256i q8s_1 = _mm256_sub_epi8(_mm256_xor_si256(s2_1, q8_1), s2_1);
+
+ aux256 = _mm256_set1_epi32(signs[2] | (signs[3] << 16));
+ aux256 = _mm256_and_si256(_mm256_shuffle_epi8(aux256,mask1), mask2);
+ const __m256i s2_2 = _mm256_cmpeq_epi8(aux256, mask2);
+ const __m256i q8s_2 = _mm256_sub_epi8(_mm256_xor_si256(s2_2, q8_2), s2_2);
+
+ signs += 4;
+
+ const __m256i dot1 = _mm256_maddubs_epi16(q2_1, q8s_1);
+ const __m256i dot2 = _mm256_maddubs_epi16(q2_2, q8s_2);
+ const uint16_t ls1 = x[i].scales[ib32/2] & 0xf;
+ const uint16_t ls2 = x[i].scales[ib32/2] >> 4;
+ const __m256i p1 = _mm256_madd_epi16(dot1, _mm256_set1_epi16(2*ls1+1));
+ const __m256i p2 = _mm256_madd_epi16(dot2, _mm256_set1_epi16(2*ls2+1));
+ sumi1 = _mm256_add_epi32(sumi1, p1);
+ sumi2 = _mm256_add_epi32(sumi2, p2);
+ }
+
+ accumf = _mm256_fmadd_ps(_mm256_set1_ps(d), _mm256_cvtepi32_ps(_mm256_add_epi32(sumi1, sumi2)), accumf);
+
+ }
+
+ *s = 0.25f * hsum_float_8(accumf);
+
+#else
+
+ float sumf = 0.f;
+ for (int i = 0; i < nb; ++i) {
+ const float d = GGML_FP16_TO_FP32(x[i].d) * y[i].d;
+ const uint8_t * restrict qs = x[i].qs;
+ const uint8_t * restrict qh = x[i].qh;
+ const uint8_t * restrict signs = x[i].signs;
+ const int8_t * restrict q8 = y[i].qs;
+ int32_t bsum = 0;
+ for (int ib32 = 0; ib32 < QK_K/32; ib32 += 2) {
+ const uint32_t ls1 = 2*(x[i].scales[ib32/2] & 0xf) + 1;
+ const uint32_t ls2 = 2*(x[i].scales[ib32/2] >> 4) + 1;
+ int32_t sumi = 0;
+ for (int l = 0; l < 4; ++l) {
+ const uint8_t * grid1 = (const uint8_t *)(iq3xs_grid + (qs[2*l+0] | ((qh[ib32+0] << (8-2*l)) & 256)));
+ const uint8_t * grid2 = (const uint8_t *)(iq3xs_grid + (qs[2*l+1] | ((qh[ib32+0] << (7-2*l)) & 256)));
+ for (int j = 0; j < 4; ++j) {
+ sumi += grid1[j] * q8[j+0] * (signs[l] & kmask_iq2xs[j+0] ? -1 : 1);
+ sumi += grid2[j] * q8[j+4] * (signs[l] & kmask_iq2xs[j+4] ? -1 : 1);
+ }
+ q8 += 8;
+ }
+ qs += 8;
+ signs += 4;
+ bsum += sumi * ls1;
+ sumi = 0;
+ for (int l = 0; l < 4; ++l) {
+ const uint8_t * grid1 = (const uint8_t *)(iq3xs_grid + (qs[2*l+0] | ((qh[ib32+1] << (8-2*l)) & 256)));
+ const uint8_t * grid2 = (const uint8_t *)(iq3xs_grid + (qs[2*l+1] | ((qh[ib32+1] << (7-2*l)) & 256)));
+ for (int j = 0; j < 4; ++j) {
+ sumi += grid1[j] * q8[j+0] * (signs[l] & kmask_iq2xs[j+0] ? -1 : 1);
+ sumi += grid2[j] * q8[j+4] * (signs[l] & kmask_iq2xs[j+4] ? -1 : 1);
+ }
+ q8 += 8;
+ }
+ qs += 8;
+ signs += 4;
+ bsum += sumi * ls2;
+ }
+ sumf += d * bsum;
+ }
+ *s = 0.25f * sumf;
+#endif
+}
+
+
#ifdef __AVX2__
static inline __m256i mul_add_epi8(const __m256i x, const __m256i y) {
const __m256i ax = _mm256_sign_epi8(x, x);
@@ -9523,6 +9831,7 @@ void ggml_vec_dot_iq4_nl_q8_0(int n, float * restrict s, size_t bs, const void *
float sumf = 0;
for (int ib = 0; ib < nb; ib += 2) {
+
q4bits.val[0] = vld1q_u8(x[ib+0].qs);
q4bits.val[1] = vld1q_u8(x[ib+1].qs);
q8b.val[0] = vld1q_s8(y[ib+0].qs);
@@ -10239,14 +10548,15 @@ typedef struct {
uint16_t * neighbours;
} iq3_entry_t;
-static iq3_entry_t iq3_data[1] = {
+static iq3_entry_t iq3_data[2] = {
+ {NULL, NULL, NULL},
{NULL, NULL, NULL},
};
static inline int iq3_data_index(int grid_size) {
(void)grid_size;
- GGML_ASSERT(grid_size == 256);
- return 0;
+ GGML_ASSERT(grid_size == 256 || grid_size == 512);
+ return grid_size == 256 ? 0 : 1;
}
static int iq3_compare_func(const void * left, const void * right) {
@@ -10278,9 +10588,44 @@ void iq3xs_init_impl(int grid_size) {
3185, 3215, 3252, 3288, 3294, 3364, 3397, 3434, 3483, 3523, 3537, 3587, 3589, 3591, 3592, 3610,
3626, 3670, 3680, 3722, 3749, 3754, 3776, 3789, 3803, 3824, 3857, 3873, 3904, 3906, 3924, 3992,
};
+ static const uint16_t kgrid_512[512] = {
+ 0, 1, 2, 5, 7, 8, 9, 10, 12, 14, 16, 17, 21, 27, 32, 34,
+ 37, 39, 41, 43, 48, 50, 57, 60, 63, 64, 65, 66, 68, 72, 73, 77,
+ 80, 83, 87, 89, 93, 100, 113, 117, 122, 128, 129, 133, 135, 136, 139, 142,
+ 145, 149, 152, 156, 162, 165, 167, 169, 171, 184, 187, 195, 201, 205, 208, 210,
+ 217, 219, 222, 228, 232, 234, 247, 249, 253, 256, 267, 271, 273, 276, 282, 288,
+ 291, 297, 312, 322, 324, 336, 338, 342, 347, 353, 357, 359, 374, 379, 390, 393,
+ 395, 409, 426, 441, 448, 450, 452, 464, 466, 470, 475, 488, 492, 512, 513, 514,
+ 516, 520, 521, 523, 525, 527, 528, 530, 537, 540, 542, 556, 558, 561, 570, 576,
+ 577, 579, 582, 584, 588, 593, 600, 603, 609, 616, 618, 632, 638, 640, 650, 653,
+ 655, 656, 660, 666, 672, 675, 685, 688, 698, 705, 708, 711, 712, 715, 721, 727,
+ 728, 732, 737, 754, 760, 771, 773, 778, 780, 793, 795, 802, 806, 808, 812, 833,
+ 840, 843, 849, 856, 858, 873, 912, 916, 919, 932, 934, 961, 963, 968, 970, 977,
+ 989, 993, 1010, 1016, 1024, 1025, 1027, 1029, 1031, 1032, 1034, 1036, 1038, 1041, 1043, 1047,
+ 1048, 1050, 1057, 1059, 1061, 1064, 1066, 1079, 1080, 1083, 1085, 1088, 1090, 1096, 1099, 1103,
+ 1106, 1109, 1113, 1116, 1122, 1129, 1153, 1156, 1159, 1169, 1171, 1176, 1183, 1185, 1195, 1199,
+ 1209, 1212, 1216, 1218, 1221, 1225, 1234, 1236, 1241, 1243, 1250, 1256, 1270, 1281, 1287, 1296,
+ 1299, 1306, 1309, 1313, 1338, 1341, 1348, 1353, 1362, 1375, 1376, 1387, 1400, 1408, 1410, 1415,
+ 1425, 1453, 1457, 1477, 1481, 1494, 1496, 1507, 1512, 1538, 1545, 1547, 1549, 1551, 1554, 1561,
+ 1563, 1565, 1570, 1572, 1575, 1577, 1587, 1593, 1601, 1603, 1605, 1612, 1617, 1619, 1632, 1648,
+ 1658, 1662, 1664, 1674, 1680, 1690, 1692, 1704, 1729, 1736, 1740, 1745, 1747, 1751, 1752, 1761,
+ 1763, 1767, 1773, 1787, 1795, 1801, 1806, 1810, 1817, 1834, 1840, 1844, 1857, 1864, 1866, 1877,
+ 1882, 1892, 1902, 1915, 1934, 1953, 1985, 1987, 2000, 2002, 2013, 2048, 2052, 2058, 2064, 2068,
+ 2071, 2074, 2081, 2088, 2104, 2114, 2119, 2121, 2123, 2130, 2136, 2141, 2147, 2153, 2157, 2177,
+ 2179, 2184, 2189, 2193, 2203, 2208, 2223, 2226, 2232, 2244, 2249, 2251, 2256, 2258, 2265, 2269,
+ 2304, 2306, 2324, 2335, 2336, 2361, 2373, 2375, 2385, 2418, 2443, 2460, 2480, 2504, 2509, 2520,
+ 2531, 2537, 2562, 2568, 2572, 2578, 2592, 2596, 2599, 2602, 2614, 2620, 2625, 2627, 2629, 2634,
+ 2641, 2650, 2682, 2688, 2697, 2707, 2712, 2718, 2731, 2754, 2759, 2760, 2775, 2788, 2793, 2805,
+ 2811, 2817, 2820, 2832, 2842, 2854, 2890, 2902, 2921, 2923, 2978, 3010, 3012, 3026, 3081, 3083,
+ 3085, 3097, 3099, 3120, 3136, 3152, 3159, 3188, 3210, 3228, 3234, 3245, 3250, 3256, 3264, 3276,
+ 3281, 3296, 3349, 3363, 3378, 3392, 3395, 3420, 3440, 3461, 3488, 3529, 3531, 3584, 3588, 3591,
+ 3600, 3602, 3614, 3616, 3628, 3634, 3650, 3657, 3668, 3683, 3685, 3713, 3716, 3720, 3726, 3729,
+ 3736, 3753, 3778, 3802, 3805, 3819, 3841, 3845, 3851, 3856, 3880, 3922, 3938, 3970, 3993, 4032,
+ };
+
const int kmap_size = 4096;
- const int nwant = 2;
- const uint16_t * kgrid = kgrid_256;
+ const int nwant = grid_size == 256 ? 2 : 3;
+ const uint16_t * kgrid = grid_size == 256 ? kgrid_256 : kgrid_512;
uint32_t * kgrid_q3xs;
int * kmap_q3xs;
uint16_t * kneighbors_q3xs;
@@ -10377,7 +10722,7 @@ void iq3xs_init_impl(int grid_size) {
}
void iq3xs_free_impl(int grid_size) {
- GGML_ASSERT(grid_size == 256);
+ GGML_ASSERT(grid_size == 256 || grid_size == 512);
const int gindex = iq3_data_index(grid_size);
if (iq3_data[gindex].grid) {
free(iq3_data[gindex].grid); iq3_data[gindex].grid = NULL;
@@ -10410,9 +10755,10 @@ static int iq3_find_best_neighbour(const uint16_t * restrict neighbours, const u
return grid_index;
}
-static void quantize_row_iq3_xxs_impl(const float * restrict x, void * restrict vy, int n, const float * restrict quant_weights) {
+static void quantize_row_iq3_xxs_impl(int grid_size, const float * restrict x, void * restrict vy, int n,
+ const float * restrict quant_weights) {
- const int gindex = iq3_data_index(256);
+ const int gindex = iq3_data_index(grid_size);
const uint32_t * kgrid_q3xs = iq3_data[gindex].grid;
const int * kmap_q3xs = iq3_data[gindex].map;
@@ -10426,9 +10772,23 @@ static void quantize_row_iq3_xxs_impl(const float * restrict x, void * restrict
const int kMaxQ = 8;
- const int nbl = n/256;
+ const int nbl = n/QK_K;
- block_iq3_xxs * y = vy;
+ ggml_fp16_t * dh;
+ uint8_t * qs;
+ int block_size;
+ if (grid_size == 256) {
+ block_iq3_xxs * y = vy;
+ dh = &y->d;
+ qs = y->qs;
+ block_size = sizeof(block_iq3_xxs);
+ } else {
+ block_iq3_s * y = vy;
+ dh = &y->d;
+ qs = y->qs;
+ block_size = sizeof(block_iq3_s);
+ }
+ int quant_size = block_size - sizeof(ggml_fp16_t);
float scales[QK_K/32];
float weight[32];
@@ -10439,20 +10799,21 @@ static void quantize_row_iq3_xxs_impl(const float * restrict x, void * restrict
bool is_on_grid[8];
bool is_on_grid_aux[8];
uint8_t block_signs[8];
- uint8_t q3[3*(QK_K/8)];
+ uint8_t q3[3*(QK_K/8)+QK_K/32];
uint32_t * scales_and_signs = (uint32_t *)(q3 + QK_K/4);
+ uint8_t * qh = q3 + 3*(QK_K/8);
for (int ibl = 0; ibl < nbl; ++ibl) {
- y[ibl].d = GGML_FP32_TO_FP16(0.f);
- memset(q3, 0, 3*QK_K/8);
+ dh[0] = GGML_FP32_TO_FP16(0.f);
+ memset(q3, 0, 3*QK_K/8+QK_K/32);
float max_scale = 0;
const float * xbl = x + QK_K*ibl;
float sumx2 = 0;
for (int i = 0; i < QK_K; ++i) sumx2 += xbl[i]*xbl[i];
- float sigma2 = sumx2/QK_K;
+ float sigma2 = 2*sumx2/QK_K;
for (int ib = 0; ib < QK_K/32; ++ib) {
const float * xb = xbl + 32*ib;
@@ -10570,7 +10931,13 @@ static void quantize_row_iq3_xxs_impl(const float * restrict x, void * restrict
printf("\n");
GGML_ASSERT(false);
}
- q3[8*ib+k] = grid_index;
+ if (grid_size == 256) {
+ q3[8*ib+k] = grid_index;
+ } else {
+ q3[8*ib+k] = grid_index & 255;
+ qh[ib] |= ((grid_index >> 8) << k);
+ }
+
}
scales_and_signs[ib] = block_signs[0] | (block_signs[1] << 7) | (block_signs[2] << 14) | (block_signs[3] << 21);
GGML_ASSERT(scale >= 0);
@@ -10579,63 +10946,25 @@ static void quantize_row_iq3_xxs_impl(const float * restrict x, void * restrict
}
if (!max_scale) {
- memset(y[ibl].qs, 0, 3*QK_K/8);
+ memset(qs, 0, quant_size);
+ dh += block_size/sizeof(ggml_fp16_t);
+ qs += block_size;
continue;
}
float d = max_scale/31;
- y[ibl].d = GGML_FP32_TO_FP16(d);
+ dh[0] = GGML_FP32_TO_FP16(d * 1.0125f); // small improvement via this fudge factor
float id = 1/d;
- float sumqx = 0, sumq2 = 0;
for (int ib = 0; ib < QK_K/32; ++ib) {
int l = nearest_int(0.5f*(id*scales[ib]-1));
l = MAX(0, MIN(15, l));
scales_and_signs[ib] |= ((uint32_t)l << 28);
- if (false) {
- const float * xb = xbl + 32*ib;
- if (quant_weights) {
- const float * qw = quant_weights + QK_K*ibl + 32*ib;
- for (int i = 0; i < 32; ++i) weight[i] = qw[i] * sqrtf(sigma2 + xb[i]*xb[i]);
- } else {
- for (int i = 0; i < 32; ++i) weight[i] = xb[i]*xb[i];
- }
- const float db = 0.25f * d * (1 + 2*l);
- for (int k = 0; k < 8; ++k) {
- const int8_t * signs = keven_signs_q2xs + 8*((scales_and_signs[ib] >> 7*(k/2)) & 127) + 4*(k%2);
- const float * xk = xb + 4*k;
- const float * wk = weight + 4*k;
- //const uint8_t * grid = (const uint8_t *)(kgrid_q3xs + q3[8*ib+k]);
- const uint8_t * grid = (const uint8_t *)(iq3xxs_grid + q3[8*ib+k]);
- float best_mse = 0; int best_index = q3[8*ib+k];
- for (int j = 0; j < 4; ++j) {
- float diff = db * grid[j] * signs[j] - xk[j];
- best_mse += wk[j] * diff * diff;
- }
- for (int idx = 0; idx < 256; ++idx) {
- //grid = (const uint8_t *)(kgrid_q3xs + idx);
- grid = (const uint8_t *)(iq3xxs_grid + idx);
- float mse = 0;
- for (int j = 0; j < 4; ++j) {
- float diff = db * grid[j] * signs[j] - xk[j];
- mse += wk[j] * diff * diff;
- }
- if (mse < best_mse) {
- best_mse = mse; best_index = idx;
- }
- }
- q3[8*ib+k] = best_index;
- //grid = (const uint8_t *)(kgrid_q3xs + best_index);
- grid = (const uint8_t *)(iq3xxs_grid + best_index);
- for (int j = 0; j < 4; ++j) {
- float q = db * grid[j] * signs[j];
- sumqx += wk[j] * q * xk[j];
- sumq2 += wk[j] * q * q;
- }
- }
- if (sumq2 > 0) y[ibl].d = GGML_FP32_TO_FP16(d*sumqx/sumq2);
- }
}
- memcpy(y[ibl].qs, q3, 3*QK_K/8);
+ memcpy(qs, q3, quant_size);
+
+ dh += block_size/sizeof(ggml_fp16_t);
+ qs += block_size;
+
}
}
@@ -10645,7 +10974,7 @@ size_t quantize_iq3_xxs(const float * src, void * dst, int nrow, int n_per_row,
int nblock = n_per_row/QK_K;
char * qrow = (char *)dst;
for (int row = 0; row < nrow; ++row) {
- quantize_row_iq3_xxs_impl(src, qrow, n_per_row, quant_weights);
+ quantize_row_iq3_xxs_impl(256, src, qrow, n_per_row, quant_weights);
src += n_per_row;
qrow += nblock*sizeof(block_iq3_xxs);
}
@@ -10660,9 +10989,226 @@ void quantize_row_iq3_xxs(const float * restrict x, void * restrict vy, int k) {
void quantize_row_iq3_xxs_reference(const float * restrict x, block_iq3_xxs * restrict y, int k) {
assert(k % QK_K == 0);
- quantize_row_iq3_xxs_impl(x, y, k, NULL);
+ quantize_row_iq3_xxs_impl(256, x, y, k, NULL);
+}
+
+static void quantize_row_iq3_s_impl(int block_size, const float * restrict x, void * restrict vy, int n,
+ const float * restrict quant_weights,
+ float * scales,
+ float * weight,
+ float * xval,
+ int8_t * L,
+ int8_t * Laux,
+ float * waux,
+ bool * is_on_grid,
+ bool * is_on_grid_aux,
+ uint8_t * block_signs) {
+
+ const int gindex = iq3_data_index(512);
+
+ const uint32_t * kgrid_q3xs = iq3_data[gindex].grid;
+ const int * kmap_q3xs = iq3_data[gindex].map;
+ const uint16_t * kneighbors_q3xs = iq3_data[gindex].neighbours;
+
+ //GGML_ASSERT(quant_weights && "missing quantization weights");
+ GGML_ASSERT(kgrid_q3xs && "forgot to call ggml_quantize_init()?");
+ GGML_ASSERT(kmap_q3xs && "forgot to call ggml_quantize_init()?");
+ GGML_ASSERT(kneighbors_q3xs && "forgot to call ggml_quantize_init()?");
+ GGML_ASSERT(n%QK_K == 0);
+
+ const int kMaxQ = 8;
+
+ const int nbl = n/QK_K;
+
+ block_iq3_s * y = vy;
+
+ const int bs4 = block_size/4;
+ const int bs8 = block_size/8;
+
+ for (int ibl = 0; ibl < nbl; ++ibl) {
+
+ memset(&y[ibl], 0, sizeof(block_iq3_s));
+ y[ibl].d = GGML_FP32_TO_FP16(0.f);
+
+ uint8_t * qs = y[ibl].qs;
+ uint8_t * qh = y[ibl].qh;
+ uint8_t * signs = y[ibl].signs;
+
+ float max_scale = 0;
+
+ const float * xbl = x + QK_K*ibl;
+ float sumx2 = 0;
+ for (int i = 0; i < QK_K; ++i) sumx2 += xbl[i]*xbl[i];
+ float sigma2 = 2*sumx2/QK_K;
+
+ for (int ib = 0; ib < QK_K/block_size; ++ib) {
+ const float * xb = xbl + block_size*ib;
+ if (quant_weights) {
+ const float * qw = quant_weights + QK_K*ibl + block_size*ib;
+ for (int i = 0; i < block_size; ++i) weight[i] = qw[i] * sqrtf(sigma2 + xb[i]*xb[i]);
+ } else {
+ for (int i = 0; i < block_size; ++i) weight[i] = xb[i]*xb[i];
+ }
+ for (int i = 0; i < block_size; ++i) waux[i] = sqrtf(weight[i]);
+ for (int k = 0; k < bs8; ++k) {
+ uint8_t s = 0;
+ for (int i = 0; i < 8; ++i) {
+ if (xb[8*k + i] >= 0) xval[8*k + i] = xb[8*k + i];
+ else {
+ xval[8*k + i] = -xb[8*k + i]; s |= (1 << i);
+ }
+ }
+ block_signs[k] = s;
+ }
+ float max = xval[0];
+ for (int i = 1; i < block_size; ++i) max = MAX(max, xval[i]);
+ if (!max) {
+ scales[ib] = 0;
+ continue;
+ }
+ float best = 0;
+ float scale = max/(2*kMaxQ-1);
+ for (int is = -15; is <= 15; ++is) {
+ float id = (2*kMaxQ-1+is*0.2f)/max;
+ float this_scale = 1/id;
+ for (int k = 0; k < bs4; ++k) {
+ for (int i = 0; i < 4; ++i) {
+ int l = nearest_int(0.5f*(id*xval[4*k+i]-1));
+ Laux[4*k+i] = MAX(0, MIN(kMaxQ-1, l));
+ }
+ uint16_t u = 0;
+ for (int i = 0; i < 4; ++i) u |= (Laux[4*k+i] << 3*i);
+ int grid_index = kmap_q3xs[u];
+ is_on_grid_aux[k] = true;
+ if (grid_index < 0) {
+ is_on_grid_aux[k] = false;
+ const uint16_t * neighbours = kneighbors_q3xs - kmap_q3xs[u] - 1;
+ grid_index = iq3_find_best_neighbour(neighbours, kgrid_q3xs, xval + 4*k, waux + 4*k, this_scale, Laux + 4*k);
+ }
+ }
+ float sumqx = 0, sumq2 = 0;
+ for (int i = 0; i < block_size; ++i) {
+ float w = weight[i];
+ float q = 2*Laux[i] + 1;
+ sumqx += w*xval[i]*q;
+ sumq2 += w*q*q;
+ }
+ if (sumq2 > 0 && sumqx*sumqx > best*sumq2) {
+ scale = sumqx/sumq2; best = scale*sumqx;
+ for (int i = 0; i < block_size; ++i) L[i] = Laux[i];
+ for (int k = 0; k < bs4; ++k) is_on_grid[k] = is_on_grid_aux[k];
+ }
+ }
+ int n_not_ongrid = 0;
+ for (int k = 0; k < bs4; ++k) if (!is_on_grid[k]) ++n_not_ongrid;
+ if (n_not_ongrid > 0 && scale > 0) {
+ float id = 1/scale;
+ for (int k = 0; k < bs4; ++k) {
+ if (is_on_grid[k]) continue;
+ uint16_t u = 0;
+ for (int i = 0; i < 4; ++i) {
+ int l = nearest_int(0.5f*(id*xval[4*k+i]-1));
+ l = MAX(0, MIN(kMaxQ-1, l));
+ u |= (l << 3*i);
+ }
+ int grid_index = kmap_q3xs[u];
+ if (grid_index < 0) {
+ const uint16_t * neighbours = kneighbors_q3xs - kmap_q3xs[u] - 1;
+ grid_index = iq3_find_best_neighbour(neighbours, kgrid_q3xs, xval + 4*k, waux + 4*k, scale, L + 4*k);
+ }
+ const int8_t * pg = (const int8_t *)(kgrid_q3xs + grid_index);
+ for (int i = 0; i < 4; ++i) L[4*k+i] = (pg[i] - 1)/2;
+ }
+ float sumqx = 0, sumq2 = 0;
+ for (int i = 0; i < block_size; ++i) {
+ float w = weight[i];
+ float q = 2*L[i] + 1;
+ sumqx += w*xval[i]*q;
+ sumq2 += w*q*q;
+ }
+ if (sumq2 > 0) scale = sumqx/sumq2;
+ }
+ if (scale < 0) {
+ // This should never happen, but just in case, flip scale so that it is positive (we use uint's to encode the scale)
+ // and correspondingly flip quant signs.
+ scale = -scale;
+ for (int k = 0; k < bs8; ++k) block_signs[k] = ~block_signs[k];
+ }
+ for (int k = 0; k < bs4; ++k) {
+ uint16_t u = 0;
+ for (int i = 0; i < 4; ++i) u |= (L[4*k+i] << 3*i);
+ int grid_index = kmap_q3xs[u];
+ if (grid_index < 0) {
+ printf("Oops: found point %u not on grid:", u);
+ for (int i = 0; i < 4; ++i) printf(" %d", L[4*k+i]);
+ printf("\n");
+ GGML_ASSERT(false);
+ }
+ qs[k] = grid_index & 255;
+ qh[(ib*bs4+k)/8] |= ((grid_index >> 8) << ((ib*bs4+k)%8));
+ }
+ qs += bs4;
+ for (int k = 0; k < bs8; ++k) signs[k] = block_signs[k];
+ signs += bs8;
+ GGML_ASSERT(scale >= 0);
+ scales[ib] = scale;
+ max_scale = MAX(max_scale, scale);
+ }
+
+ if (!max_scale) {
+ continue;
+ }
+
+ float d = max_scale/31;
+ y[ibl].d = GGML_FP32_TO_FP16(d);
+ float id = 1/d;
+ for (int ib = 0; ib < QK_K/block_size; ib += 2) {
+ int l1 = nearest_int(0.5f*(id*scales[ib+0]-1));
+ l1 = MAX(0, MIN(15, l1));
+ int l2 = nearest_int(0.5f*(id*scales[ib+1]-1));
+ l2 = MAX(0, MIN(15, l2));
+ y[ibl].scales[ib/2] = l1 | (l2 << 4);
+ }
+
+ }
+}
+
+#define IQ3S_BLOCK_SIZE 32
+size_t quantize_iq3_s(const float * src, void * dst, int nrow, int n_per_row, int64_t * hist, const float * quant_weights) {
+ (void)hist;
+ GGML_ASSERT(n_per_row%QK_K == 0);
+ int nblock = n_per_row/QK_K;
+ float scales[QK_K/IQ3S_BLOCK_SIZE];
+ float weight[IQ3S_BLOCK_SIZE];
+ float xval[IQ3S_BLOCK_SIZE];
+ int8_t L[IQ3S_BLOCK_SIZE];
+ int8_t Laux[IQ3S_BLOCK_SIZE];
+ float waux[IQ3S_BLOCK_SIZE];
+ bool is_on_grid[IQ3S_BLOCK_SIZE/4];
+ bool is_on_grid_aux[IQ3S_BLOCK_SIZE/4];
+ uint8_t block_signs[IQ3S_BLOCK_SIZE/8];
+ char * qrow = (char *)dst;
+ for (int row = 0; row < nrow; ++row) {
+ quantize_row_iq3_s_impl(IQ3S_BLOCK_SIZE, src, qrow, n_per_row, quant_weights,
+ scales, weight, xval, L, Laux, waux, is_on_grid, is_on_grid_aux, block_signs);
+ src += n_per_row;
+ qrow += nblock*sizeof(block_iq3_s);
+ }
+ return nrow * nblock * sizeof(block_iq3_s);
+}
+
+void quantize_row_iq3_s(const float * restrict x, void * restrict vy, int k) {
+ assert(k % QK_K == 0);
+ block_iq3_s * restrict y = vy;
+ quantize_row_iq3_s_reference(x, y, k);
}
+void quantize_row_iq3_s_reference(const float * restrict x, block_iq3_s * restrict y, int k) {
+ assert(k % QK_K == 0);
+ quantize_iq3_s(x, y, 1, k, NULL, NULL);
+}
+
+
// =================================== 1.5 bpw ===================================================
static int iq1_find_best_neighbour(const uint16_t * restrict neighbours, const uint64_t * restrict grid,
generated by cgit v1.2.3 (git 2.25.1) at 2025-08-18 07:17:37 +0000