From dfcc8a9cf3820214a2b1348b33a257509c688bef Mon Sep 17 00:00:00 2001
From: Iwan Kawrakow <iwan.kawrakow@gmail.com>
Date: Sat, 9 Nov 2024 09:32:00 +0200
Subject: [PATCH] iq3_kt WIP: slowly improving

PPL(LLaMA-3.1-8B-Instruct, 8192) is now 6.7892
---
 ggml/src/ggml-cuda/convert.cu |  2 +-
 ggml/src/iqk/iqk_quantize.cpp | 12 ++++++------
 2 files changed, 7 insertions(+), 7 deletions(-)

diff --git a/ggml/src/ggml-cuda/convert.cu b/ggml/src/ggml-cuda/convert.cu
index 47fc0dbc..43f07465 100644
--- a/ggml/src/ggml-cuda/convert.cu
+++ b/ggml/src/ggml-cuda/convert.cu
@@ -391,7 +391,7 @@ static __global__ void dequantize_block_iq3_kt(const void * __restrict__ vx, dst
     dst_t * y = yy + ii*QK_K + 8*ib;
     uint32_t idx1 = x[i].ql[2*ib+0] + ((x[i].qh[(2*ib+0)%32] << (8-4*((2*ib+0)/32))) & 0xf00) + 4096;
     uint32_t idx2 = x[i].ql[2*ib+1] + ((x[i].qh[(2*ib+1)%32] << (8-4*((2*ib+1)/32))) & 0xf00) + 4096;
-    const float dl = scale * iq4k_values[((x[i].scales[(ib/4)%4] >> 4*(ib/16)) & 0xf)] * 31.75f * 1.015f;
+    const float dl = scale * iq4k_values[((x[i].scales[(ib/4)%4] >> 4*(ib/16)) & 0xf)+16] * 31.75f * 1.015f;
     uint32_t s[2];
     const half * h = (const half *)s;
     for (int j = 0; j < 4; ++j) {
diff --git a/ggml/src/iqk/iqk_quantize.cpp b/ggml/src/iqk/iqk_quantize.cpp
index 641f93c3..38dbc54b 100644
--- a/ggml/src/iqk/iqk_quantize.cpp
+++ b/ggml/src/iqk/iqk_quantize.cpp
@@ -4000,13 +4000,13 @@ void quantize_row_iq3_kt_impl(const float * x, void * vy, int n_per_row, const f
 
     }
 
-    float d = max_scale/iq4k_values[0];
+    float d = max_scale/iq4k_values[16];
     float id = d ? 1/d : 0.f;
     for (int ibl = 0; ibl < nblock; ++ibl) {
         auto scales = all_scales + ibl*Q::kNblock;
         for (int ib = 0; ib < Q::kNblock/2; ++ib) {
-            int ls1 = best_index_iq4nl(iq4k_values, id*scales[ib]);
-            int ls2 = best_index_iq4nl(iq4k_values, id*scales[ib + Q::kNblock/2]);
+            int ls1 = best_index_iq4nl(iq4k_values+16, id*scales[ib]);
+            int ls2 = best_index_iq4nl(iq4k_values+16, id*scales[ib + Q::kNblock/2]);
             y[ibl].scales[ib] = ls1 | (ls2 << 4);
         }
     }
@@ -4035,7 +4035,7 @@ void quantize_row_iq3_kt_impl(const float * x, void * vy, int n_per_row, const f
                 } else {
                     for (int j = 0; j < Q::kBlockSize; ++j) weight[j] = 0.25f*sigma2 + xb[j]*xb[j];
                 }
-                int ls = iq4k_values[(y[ibl].scales[ib%(Q::kNblock/2)] >> 4*(ib/(Q::kNblock/2))) & 0xf];
+                int ls = iq4k_values[((y[ibl].scales[ib%(Q::kNblock/2)] >> 4*(ib/(Q::kNblock/2))) & 0xf)+16];
                 float dl = d*ls;
                 quantizer.find_best_match(dl, xb, weight, best_idx);
 
@@ -4104,8 +4104,8 @@ void dequantize_row_iq3_kt(const block_iq3_kt * x, float * y, int64_t k) {
         auto qlh = qll + kNumGroups/2;
         int jj = 0;
         for (int ib = 0; ib < Q::kNblock/2; ++ib) {
-            float sl = d * iq4k_values[x[ibl].scales[ib] & 0xf];
-            float sh = d * iq4k_values[x[ibl].scales[ib] >>  4];
+            float sl = d * iq4k_values[(x[ibl].scales[ib] & 0xf)+16];
+            float sh = d * iq4k_values[(x[ibl].scales[ib] >>  4)+16];
             for (int ig = 0; ig < Q::kNg; ++ig) {
                 uint16_t ul = qll[jj] | ((x[ibl].qh[jj] << 8) & 0xf00);
                 uint16_t uh = qlh[jj] | ((x[ibl].qh[jj] << 4) & 0xf00);