From e558992f0ce34615bbc74c7ffc00fbaf617dba34 Mon Sep 17 00:00:00 2001 From: Iwan Kawrakow Date: Sat, 7 Jun 2025 12:30:37 +0300 Subject: [PATCH] New iq4_kt trellis The new trellis generates int8_t values via sum_as_uint8_t[(ka * idx + kb) & 0x3f33f3f3f] - 126. CUDA dequantize works. AVX2 case Ny > 32 works, and we get 273 t/s for L3-8B. PPL is on par or even slightly lower than original QTIP trellis. --- ggml/src/ggml-cuda/convert.cu | 12 +-- ggml/src/ggml.c | 11 +- ggml/src/iqk/iqk_gemm_ktquants.cpp | 157 +++++++++++++++++++++++++++-- ggml/src/iqk/iqk_mul_mat.cpp | 6 +- ggml/src/iqk/iqk_quantize.cpp | 29 ++++-- 5 files changed, 181 insertions(+), 34 deletions(-) diff --git a/ggml/src/ggml-cuda/convert.cu b/ggml/src/ggml-cuda/convert.cu index 01b7250e..b1c447ab 100644 --- a/ggml/src/ggml-cuda/convert.cu +++ b/ggml/src/ggml-cuda/convert.cu @@ -343,13 +343,9 @@ inline __device__ int nearest_int(float fval) { float __device__ __forceinline__ trellis_next(uint32_t& val) { constexpr uint32_t ka = 89226354; constexpr uint32_t kb = 64248484; - constexpr uint32_t kmask = 0x8fff8fff; - constexpr uint32_t km32 = 0x3b603b60; - uint32_t s; - const half * h = (const half *)&s; val = ka*val + kb; - s = (val & kmask) ^ km32; - return (float)(h[0]+h[1]); + //return ggml_cuda_dp4a(val & 0x3f3f3f3f, 0x01010101, 0x82828282); + return ggml_cuda_dp4a(val & 0x3f3f3f3f, 0x01010101, -126); } template @@ -367,7 +363,7 @@ static __global__ void dequantize_block_iq2_kt(const void * __restrict__ vx, dst dst_t * y = yy + ii*QK_K + 8*ib; const uint16_t * ql = (const uint16_t *)x[i].ql; uint32_t idx = ql[ib] + 4096; - const float dl = scale * iq4k_values[((x[i].scales[(ib/4)%4] >> 4*(ib/16)) & 0xf)] * 31.75f * 1.05f; + const float dl = scale * iq4k_values[((x[i].scales[(ib/4)%4] >> 4*(ib/16)) & 0xf)] * 1.05f; for (int j = 0; j < 8; ++j) { y[j] = dl * trellis_next(idx); } @@ -401,7 +397,7 @@ static __global__ void dequantize_block_iq4_kt(const void * __restrict__ vx, dst int64_t ii = blockIdx.x; int64_t row = (QK_K * ii) / n_per_row; const float * dptr = (const float *)((const char *)vx + row * row_size); - float scale = dptr[0] * 31.75f * 1.01f; + float scale = dptr[0] * 1.00f; float row_av = dptr[1]; const block_iq4_kt * x = (const block_iq4_kt *)(dptr + 2); const int64_t i = ii - (row*n_per_row)/QK_K; diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index 69b1b46d..b1b9b57b 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -1630,11 +1630,16 @@ static const ggml_type_traits_t type_traits[GGML_TYPE_COUNT] = { .from_float = quantize_row_iq4_kt, .from_float_ref = (ggml_from_float_t)quantize_row_iq4_kt_ref, .vec_dot = vec_dot_iq4_kt_q8_k, -#ifdef __ARM_NEON - .vec_dot_type = GGML_TYPE_F16, +#if defined __AVX2__ + .vec_dot_type = GGML_TYPE_Q8_2_X4, #else - .vec_dot_type = GGML_TYPE_F32, + .vec_dot_type = GGML_TYPE_Q8_0_X4, #endif +//#ifdef __ARM_NEON +// .vec_dot_type = GGML_TYPE_F16, +//#else +// .vec_dot_type = GGML_TYPE_F32, +//#endif .nrows = 1, .row_meta_size = 8, }, diff --git a/ggml/src/iqk/iqk_gemm_ktquants.cpp b/ggml/src/iqk/iqk_gemm_ktquants.cpp index bc7bcf8b..0a8d2d03 100644 --- a/ggml/src/iqk/iqk_gemm_ktquants.cpp +++ b/ggml/src/iqk/iqk_gemm_ktquants.cpp @@ -97,6 +97,43 @@ struct Trellis2 { } }; + +struct Trellis3 { + constexpr static uint32_t ka = 89226354; + constexpr static uint32_t kb = 64248484; + constexpr static uint32_t ka1 = ka*ka; + constexpr static uint32_t kb1 = kb*ka+kb; + constexpr static uint32_t ka2 = ka1*ka; + constexpr static uint32_t kb2 = kb1*ka+kb; + constexpr static uint32_t ka3 = ka2*ka; + constexpr static uint32_t kb3 = kb2*ka+kb; + const __m256i mka = _mm256_setr_epi32(ka, ka1, ka2, ka3, ka, ka1, ka2, ka3); + const __m256i mkb = _mm256_setr_epi32(kb, kb1, kb2, kb3, kb, kb1, kb2, kb3); + const __m256i shuffle = _mm256_set_epi32(7, 3, 6, 2, 5, 1, 4, 0); + + inline __m256i next8(uint32_t val1, uint32_t val2) const { + __m256i mval = MM256_SET_M128I(_mm_set1_epi32(val2), _mm_set1_epi32(val1)); + return _mm256_add_epi32(_mm256_mullo_epi32(mval, mka), mkb); + } + inline __m256 gen8(uint32_t val1, uint32_t val2) const { + auto v8 = _mm256_and_si256(next8(val1, val2), _mm256_set1_epi32(0x3f3f3f3f)); + auto i8 = _mm256_dpbusd_epi32(_mm256_set1_epi32(-126), _mm256_set1_epi32(0x01010101), v8); + return _mm256_cvtepi32_ps(i8); + } + inline __m256i next32(const uint32_t * val) const { + __m256i aux[4]; + for (int i = 0; i < 4; ++i) { + auto i8 = _mm256_and_si256(next8(val[2*i+0], val[2*i+1]), _mm256_set1_epi32(0x3f3f3f3f)); + aux[i] = _mm256_dpbusd_epi32(_mm256_set1_epi32(-126), _mm256_set1_epi32(0x01010101), i8); + } + aux[0] = _mm256_packs_epi32(aux[0], aux[1]); // 0, 1, 2, 3, 8, 9, 10, 11, 4, 5, 6, 7, 12, 13, 14, 15 + aux[2] = _mm256_packs_epi32(aux[2], aux[3]); // 16, 17, 18, 19, 24, 25, 26, 27, 20, 21, 22, 23, 28, 29, 30, 31 + aux[0] = _mm256_packs_epi16(aux[0], aux[2]); // 0, 1, 2, 3, 8, 9, 10, 11, 16, 17, 18, 19, 24, 25, 26, 27 + // 4, 5, 6, 7, 12, 13, 14, 15, 20, 21, 22, 23, 28, 29, 30, 31 + return _mm256_permutevar8x32_epi32(aux[0], shuffle); + } +}; + void iqk_dequantize_iq2_kt(int n, const void * vx, size_t bx, float * y, size_t stride_y, int nrc_x) { GGML_ASSERT(n%QK_K == 0); const int nb = n/QK_K; @@ -315,19 +352,121 @@ void mul_mat_iq3_kt_F32_T(int n, const void * vx, size_t bx, const DataInfo& inf } } +// Q8_0 repacking: +// for (int ib = 0; ib < nblock; ++ib) { +// for (int k = 0; k < 8; ++k) y[ib].d[k] = x8[k][ib].d; +// for (int l = 0; l < 4; ++l) { +// for (int k = 0; k < 8; ++k) for (int i = 0; i < 4; ++i) { +// y[ib].qs[32*l+4*k+i+ 0] = x8[k][ib].qs[i+4*l+ 0]; +// y[ib].qs[32*l+4*k+i+128] = x8[k][ib].qs[i+4*l+16]; +// as uint32_t +// y[ib].qs[8*l+k+ 0] = x8[k][ib].qs[l+ 0]; +// y[ib].qs[8*l+k+32] = x8[k][ib].qs[l+16]; +// } +// } +// } + +void iqk_dequantize_iq4_kt_q80_r8(int n, const void * vx, size_t bx, void * vy, int nrc_x) { + GGML_ASSERT(n%QK_K == 0); + GGML_ASSERT(nrc_x%8 == 0); + const int nb = n/QK_K; + constexpr int kNumGroups = 64; + + Trellis3 trellis; + + block_q8_0_r8 * y = (block_q8_0_r8 *)vy; + + const block_iq4_kt * x8[8]; + float dkt[8]; + int32_t ls[8]; + uint32_t idx0[8], idx[16]; + + for (int ix = 0; ix < nrc_x; ix += 8) { + for (int k = 0; k < 8; ++k) { + const float * dptr = (const float *)((const char*)vx + (ix+k)*bx); + dkt[k] = dptr[0]; + x8[k] = (const block_iq4_kt *)(dptr + 2); + } + auto vd = _mm256_loadu_ps(dkt); + + for (int i = 0; i < nb; ++i) { + for (int ib = 0; ib < QK_K/32; ++ib) { + for (int k = 0; k < 8; ++k) { + ls[k] = ((x8[k][i].qs[ib] & 0xff) >> 1) - 64; + idx0[k] = ((x8[k][i].qs[ib] & 1) << 15) + 4096; + } + auto scales = _mm256_mul_ps(vd, _mm256_cvtepi32_ps(_mm256_loadu_si256((const __m256i *)ls))); + _mm_storeu_si128((__m128i *)y[ib].d, _mm256_cvtps_ph(scales, _MM_FROUND_TO_NEAREST_INT)); + //for (int k = 0; k < 8; ++k) { + // auto shb = x8[k][i].qs; + // const uint8_t * ql = (const uint8_t *)(shb + 8); + // const uint8_t * qh = ql + kNumGroups; + // for (int ib = 0; ib < 4; ++ib) { + // uint32_t offset1 = ((shb[ib+0] & 1) << 15) + 4096; + // uint32_t offset2 = ((shb[ib+4] & 1) << 15) + 4096; + // for (int j = 0; j < 4; ++j) { + // const uint32_t sh1 = shb[ib+0] >> (8 + 6*j); + // const uint32_t sh2 = shb[ib+4] >> (8 + 6*j); + // idx[64*ib + 16*j + k ] = ql[8*ib+2*j+ 0] + ((qh[8*ib+2*j+0] << 8) & 0xf00) + ((sh1 & 7) << 12) + offset1; + // idx[64*ib + 16*j + k + 8] = ql[8*ib+2*j+ 1] + ((qh[8*ib+2*j+1] << 8) & 0xf00) + ((sh1 & 56) << 9) + offset1; + // idx[64*ib + 16*j + k + 256] = ql[8*ib+2*j+32] + ((qh[8*ib+2*j+0] << 4) & 0xf00) + ((sh2 & 7) << 12) + offset2; + // idx[64*ib + 16*j + k + 264] = ql[8*ib+2*j+33] + ((qh[8*ib+2*j+1] << 4) & 0xf00) + ((sh2 & 56) << 9) + offset2; + // //uint32_t val1 = ql[8*ib+2*j+ 0] + ((qh[8*ib+2*j+0] << 8) & 0xf00) + ((sh1 & 7) << 12) + offset1; + // //uint32_t val2 = ql[8*ib+2*j+32] + ((qh[8*ib+2*j+0] << 4) & 0xf00) + ((sh2 & 7) << 12) + offset2; + // //uint32_t val3 = ql[8*ib+2*j+ 1] + ((qh[8*ib+2*j+1] << 8) & 0xf00) + ((sh1 & 56) << 9) + offset1; + // //uint32_t val4 = ql[8*ib+2*j+33] + ((qh[8*ib+2*j+1] << 4) & 0xf00) + ((sh2 & 56) << 9) + offset2; + // //auto x_val1 = _mm256_fmadd_ps(scale1, trellis.gen8(val1, val3), dav); + // //auto x_val2 = _mm256_fmadd_ps(scale2, trellis.gen8(val2, val4), dav); + // //_mm256_storeu_ps(y + i*QK_K + 32*ib + 8*j, x_val1); + // //_mm256_storeu_ps(y + i*QK_K + 32*ib + 8*j + QK_K/2, x_val2); + // } + // } + //} + //for (int j = 0; j < 64; ++j) { + // _mm256_storeu_si256((__m256i *)y[j/8].qs+(j%8), trellis.next32(idx+8*j)); + //} + //int shift1 = 8 - 4*(ib/4); + //for (int j = 0; j < 4; ++j) { + // for (int k = 0; k < 8; ++k) { + // const uint8_t * ql = (const uint8_t *)(x8[k][i].qs + 8); + // const uint8_t * qh = ql + kNumGroups; + // const uint32_t sh = x8[k][i].qs[ib] >> (8 + 6*j); + // idx[k+0] = ql[8*ib+2*j+0] + ((qh[8*(ib%4)+2*j+0] << shift1) & 0xf00) + ((sh & 7) << 12) + idx0[k]; + // idx[k+8] = ql[8*ib+2*j+1] + ((qh[8*(ib%4)+2*j+1] << shift1) & 0xf00) + ((sh & 56) << 9) + idx0[k]; + // } + // _mm256_storeu_si256((__m256i *)y[ib].qs+2*j+0, trellis.next32(idx+0)); + // _mm256_storeu_si256((__m256i *)y[ib].qs+2*j+1, trellis.next32(idx+8)); + //} + int shift1 = 8 - 4*(ib/4); + for (int j = 0; j < 8; ++j) { + for (int k = 0; k < 8; ++k) { + const uint8_t * ql = (const uint8_t *)(x8[k][i].qs + 8); + const uint8_t * qh = ql + kNumGroups; + const uint32_t sh = x8[k][i].qs[ib] >> (8 + 3*j); + idx[k+0] = ql[8*ib+j] + ((qh[8*(ib%4)+j] << shift1) & 0xf00) + ((sh & 7) << 12) + idx0[k]; + } + _mm256_storeu_si256((__m256i *)y[ib].qs+j, trellis.next32(idx)); + } + } + y += 8; // = QK_K/32; + } + + } +} + void iqk_dequantize_iq4_kt(int n, const void * vx, size_t bx, float * y, size_t stride_y, int nrc_x) { GGML_ASSERT(n%QK_K == 0); const int nb = n/QK_K; constexpr int kNumGroups = 64; - Trellis2 trellis; + Trellis3 trellis; union { __m256 vec; float val[8]; } s_helper; union { __m256i vec; uint32_t val[8]; } o_helper; for (int ix = 0; ix < nrc_x; ++ix) { const float * dptr = (const float *)((const char*)vx + ix*bx); - auto d = _mm256_set1_ps(dptr[0] * 31.75f * 1.01f); + auto d = _mm256_set1_ps(dptr[0]); auto dav = _mm256_set1_ps(dptr[1]); const block_iq4_kt * x = (const block_iq4_kt *)(dptr + 2); @@ -349,8 +488,8 @@ void iqk_dequantize_iq4_kt(int n, const void * vx, size_t bx, float * y, size_t uint32_t val2 = ql[8*ib+2*j+32] + ((qh[8*ib+2*j+0] << 4) & 0xf00) + ((sh2 & 7) << 12) + o_helper.val[ib+4]; uint32_t val3 = ql[8*ib+2*j+ 1] + ((qh[8*ib+2*j+1] << 8) & 0xf00) + ((sh1 & 56) << 9) + o_helper.val[ib+0]; uint32_t val4 = ql[8*ib+2*j+33] + ((qh[8*ib+2*j+1] << 4) & 0xf00) + ((sh2 & 56) << 9) + o_helper.val[ib+4]; - auto x_val1 = _mm256_fmadd_ps(scale1, trellis_gen8(trellis.next8(val1, val3)), dav); - auto x_val2 = _mm256_fmadd_ps(scale2, trellis_gen8(trellis.next8(val2, val4)), dav); + auto x_val1 = _mm256_fmadd_ps(scale1, trellis.gen8(val1, val3), dav); + auto x_val2 = _mm256_fmadd_ps(scale2, trellis.gen8(val2, val4), dav); _mm256_storeu_ps(y + i*QK_K + 32*ib + 8*j, x_val1); _mm256_storeu_ps(y + i*QK_K + 32*ib + 8*j + QK_K/2, x_val2); @@ -370,7 +509,7 @@ void mul_mat_iq4_kt_F32_T(int n, const void * vx, size_t bx, const DataInfo& inf const int nb = n/QK_K; constexpr int kNumGroups = 64; - Trellis2 trellis; + Trellis3 trellis; union { __m256 vec; float val[8]; } s_helper; union { __m256i vec; uint32_t val[8]; } o_helper; @@ -389,7 +528,7 @@ void mul_mat_iq4_kt_F32_T(int n, const void * vx, size_t bx, const DataInfo& inf for (int ix = 0; ix < nrc_x; ++ix) { const float * dptr = (const float *)((const char*)vx + ix*bx); - auto d = _mm256_set1_ps(dptr[0] * 31.75f * 1.01f); + auto d = _mm256_set1_ps(dptr[0]); auto dav = dptr[1]; const block_iq4_kt * x = (const block_iq4_kt *)(dptr + 2); @@ -413,8 +552,8 @@ void mul_mat_iq4_kt_F32_T(int n, const void * vx, size_t bx, const DataInfo& inf uint32_t val2 = ql[8*ib+2*j+32] + ((qh[8*ib+2*j+0] << 4) & 0xf00) + ((sh2 & 7) << 12) + o_helper.val[ib+4]; uint32_t val3 = ql[8*ib+2*j+ 1] + ((qh[8*ib+2*j+1] << 8) & 0xf00) + ((sh1 & 56) << 9) + o_helper.val[ib+0]; uint32_t val4 = ql[8*ib+2*j+33] + ((qh[8*ib+2*j+1] << 4) & 0xf00) + ((sh2 & 56) << 9) + o_helper.val[ib+4]; - auto x_val1 = _mm256_mul_ps(scale1, trellis_gen8(trellis.next8(val1, val3))); - auto x_val2 = _mm256_mul_ps(scale2, trellis_gen8(trellis.next8(val2, val4))); + auto x_val1 = _mm256_mul_ps(scale1, trellis.gen8(val1, val3)); + auto x_val2 = _mm256_mul_ps(scale2, trellis.gen8(val2, val4)); if constexpr (nrc_y == 1) { auto y1 = _mm256_load_ps(y[0] + i*QK_K+32*ib+8*j+ 0); auto y2 = _mm256_load_ps(y[0] + i*QK_K+32*ib+8*j+128); @@ -474,7 +613,7 @@ bool iqk_dequantize_ktquants(int type, int n, const void * vx, size_t bx, void * switch (type) { case GGML_TYPE_IQ2_KT: iqk_dequantize_iq2_kt(n, vx, bx, (float *)y, stride_y, nrc_x); break; case GGML_TYPE_IQ3_KT: iqk_dequantize_iq3_kt(n, vx, bx, (float *)y, stride_y, nrc_x); break; - case GGML_TYPE_IQ4_KT: iqk_dequantize_iq4_kt(n, vx, bx, (float *)y, stride_y, nrc_x); break; + case GGML_TYPE_IQ4_KT: iqk_dequantize_iq4_kt_q80_r8(n, vx, bx, y, nrc_x); break; default: return false; } return true; diff --git a/ggml/src/iqk/iqk_mul_mat.cpp b/ggml/src/iqk/iqk_mul_mat.cpp index 6925e6a6..ce67159c 100644 --- a/ggml/src/iqk/iqk_mul_mat.cpp +++ b/ggml/src/iqk/iqk_mul_mat.cpp @@ -236,9 +236,6 @@ struct MulMat { static inline ggml_type is_dequant_better(ggml_type type, int nrc_y) { #ifdef __AVX2__ switch (type) { - case GGML_TYPE_IQ2_KT : return nrc_y >= 32 ? GGML_TYPE_F32 : type; - case GGML_TYPE_IQ3_KT : return nrc_y >= 32 ? GGML_TYPE_F32 : type; - case GGML_TYPE_IQ4_KT : return nrc_y >= 32 ? GGML_TYPE_F32 : type; case GGML_TYPE_IQ2_XXS: return nrc_y >= 32 ? GGML_TYPE_Q8_K_R8 : type; case GGML_TYPE_IQ2_XS : return nrc_y >= 32 ? GGML_TYPE_Q8_K_R8 : type; case GGML_TYPE_IQ2_S : return nrc_y >= 16 ? GGML_TYPE_Q8_K_R8 : type; @@ -267,6 +264,9 @@ struct MulMat { case GGML_TYPE_Q6_0 : return nrc_y >= 32 ? GGML_TYPE_Q8_0_R8 : type; case GGML_TYPE_IQ4_NL : return nrc_y >= 32 ? GGML_TYPE_Q8_0_R8 : type; case GGML_TYPE_Q8_0 : return nrc_y >= 32 ? GGML_TYPE_Q8_0_R8 : type; + case GGML_TYPE_IQ2_KT : return nrc_y >= 32 ? GGML_TYPE_F32 : type; + case GGML_TYPE_IQ3_KT : return nrc_y >= 32 ? GGML_TYPE_F32 : type; + case GGML_TYPE_IQ4_KT : return nrc_y >= 32 ? GGML_TYPE_Q8_0_R8 : type; default: break; } #else diff --git a/ggml/src/iqk/iqk_quantize.cpp b/ggml/src/iqk/iqk_quantize.cpp index abd4be61..7061c8bf 100644 --- a/ggml/src/iqk/iqk_quantize.cpp +++ b/ggml/src/iqk/iqk_quantize.cpp @@ -7408,7 +7408,7 @@ public: constexpr static int kNg = kBlockSize/kGroupSize; constexpr static int kNblock = kSuperBlockSize/kBlockSize; constexpr static int kNumVal = 1 << num_bits; // i.e, 16 bits per group of 8 - constexpr static float kScale = 31.75f; + constexpr static float kScale = 1.f; //31.75f; constexpr static bool kVerbose = false; QuantizerIQKT(int num_clusters, int num_neighbours, int offset = 4096); @@ -7421,15 +7421,19 @@ public: static inline void set_values(uint32_t i, float * result, float scale, int offset = 4096) { constexpr uint32_t ka = 89226354; constexpr uint32_t kb = 64248484; - constexpr uint32_t kmask = 0x8fff8fff; - constexpr uint32_t km32 = 0x3b603b60; + //constexpr uint32_t kmask = 0x8fff8fff; + //constexpr uint32_t km32 = 0x3b603b60; uint32_t x = i + offset; + uint32_t s; + auto i8 = (const int8_t *)&s; for (int k = 0; k < kGroupSize; ++k) { x = ka*x + kb; - uint32_t s = (x & kmask) ^ km32; - float val = GGML_FP16_TO_FP32(s & 65535) + GGML_FP16_TO_FP32(s >> 16); - if constexpr (is_abs) result[k] = scale*std::abs(val); - else result[k] = scale*val; + s = x & 0x3f3f3f3f; + result[k] = scale*(i8[0] + i8[1] + i8[2] + i8[3] - 126.f); + //uint32_t s = (x & kmask) ^ km32; + //float val = GGML_FP16_TO_FP32(s & 65535) + GGML_FP16_TO_FP32(s >> 16); + //if constexpr (is_abs) result[k] = scale*std::abs(val); + //else result[k] = scale*val; } } @@ -8209,7 +8213,7 @@ size_t quantize_iq2_kt(const float * src, void * dst, int64_t nrows, int64_t n_p void dequantize_row_iq2_kt(const block_iq2_kt * x, float * y, int64_t k) { assert(k % QuantizerIQ2KT::kSuperBlockSize == 0); #ifdef __AVX2__ - if (iqk_dequantize_ktquants(GGML_TYPE_IQ2_KT, k, x, 0, y, 0, 1)) return; + //if (iqk_dequantize_ktquants(GGML_TYPE_IQ2_KT, k, x, 0, y, 0, 1)) return; #endif const int nb = k / QuantizerIQ2KT::kSuperBlockSize; const float * dptr = (const float *)x; @@ -8560,7 +8564,10 @@ void quantize_row_iq4_kt_impl(const float * x, void * vy, int n_per_row, const f row_av += x[j]; amax_row = std::max(amax_row, std::abs(x[j])); } - row_av /= n_per_row; + //!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! + //row_av /= n_per_row; + row_av = 0; + //!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! dptr[1] = row_av; if (!amax_row) { dptr[0] = 0.f; @@ -8593,7 +8600,7 @@ void quantize_row_iq4_kt_impl(const float * x, void * vy, int n_per_row, const f continue; } float best = 0; - float scale_0 = std::max(92.f, 127.f*amax/amax_row); + float scale_0 = std::max(90.f, 124.f*amax/amax_row); for (int itry = -kNtry; itry <= kNtry; ++itry) { quantizer1.find_best_match( amax/(8.f*itry + scale_0), xaux, weight, best_idx); auto [dp, score_p] = quantizer1.find_best_scale(xaux, weight, best_idx); @@ -8724,7 +8731,7 @@ size_t quantize_iq4_kt(const float * src, void * dst, int64_t nrows, int64_t n_p void dequantize_row_iq4_kt(const block_iq4_kt * x, float * y, int64_t k) { #ifdef __AVX2__ - if (iqk_dequantize_ktquants(GGML_TYPE_IQ4_KT, k, x, 0, y, 0, 1)) return; + //if (iqk_dequantize_ktquants(GGML_TYPE_IQ4_KT, k, x, 0, y, 0, 1)) return; #endif using Q = QuantizerIQ4KT; assert(k % Q::kSuperBlockSize == 0);