From a08501ee5216402458d3d3e9b9af5763705eaffe Mon Sep 17 00:00:00 2001 From: Kawrakow Date: Thu, 6 Feb 2025 18:45:28 +0200 Subject: Rename q4_0_r4, q8_0_r4 and iq4_xs_r4 to _r8 (#189) * Rename q4_0_r4 to q4_0_r8 to reflect actual row interleaving * Rename q8_0_r4 to q8_0_r8 to reflect actual row interleaving * Rename iq4_xs_r4 to iq4_xs_r8 to reflect actual row interleaving --------- Co-authored-by: Iwan Kawrakow --- ggml/src/ggml-common.h | 4 +- ggml/src/ggml-quants.c | 6 +-- ggml/src/ggml.c | 90 +++++++++++++++---------------- ggml/src/iqk/iqk_mul_mat.cpp | 120 +++++++++++++++++++++--------------------- ggml/src/iqk/iqk_quantize.cpp | 72 ++++++++++++------------- ggml/src/iqk/iqk_quantize.h | 32 +++++------ 6 files changed, 162 insertions(+), 162 deletions(-) (limited to 'ggml/src') diff --git a/ggml/src/ggml-common.h b/ggml/src/ggml-common.h index 679353be..0d014c23 100644 --- a/ggml/src/ggml-common.h +++ b/ggml/src/ggml-common.h @@ -562,8 +562,8 @@ typedef struct { uint8_t scales_h[QK_K/16]; uint8_t scales_l[QK_K/ 8]; uint8_t qs[QK_K*4]; -} block_iq4_xs_r4; -static_assert(sizeof(block_iq4_xs_r4) == 8*sizeof(block_iq4_xs), "wrong iq4_xs_rs block size/padding"); +} block_iq4_xs_r8; +static_assert(sizeof(block_iq4_xs_r8) == 8*sizeof(block_iq4_xs), "wrong iq4_xs_rs block size/padding"); typedef struct { uint8_t scales[QK_K/32]; diff --git a/ggml/src/ggml-quants.c b/ggml/src/ggml-quants.c index d32a583f..fe7de167 100644 --- a/ggml/src/ggml-quants.c +++ b/ggml/src/ggml-quants.c @@ -15193,7 +15193,7 @@ bool ggml_validate_row_data(enum ggml_type type, const void * data, size_t nbyte case GGML_TYPE_IQ4_KS: break; case GGML_TYPE_IQ4_KSS: break; case GGML_TYPE_IQ4_NL_R4: break; - case GGML_TYPE_IQ4_XS_R4: break; + case GGML_TYPE_IQ4_XS_R8: break; case GGML_TYPE_IQ2_XXS_R4: break; case GGML_TYPE_IQ2_XS_R4: break; case GGML_TYPE_IQ3_XXS_R4: break; @@ -15201,10 +15201,10 @@ bool ggml_validate_row_data(enum ggml_type type, const void * data, size_t nbyte case GGML_TYPE_IQ2_S_R4: break; case GGML_TYPE_IQ1_S_R4: break; case GGML_TYPE_IQ1_M_R4: break; - case GGML_TYPE_Q4_0_R4: break; + case GGML_TYPE_Q4_0_R8: break; case GGML_TYPE_Q5_0_R4: break; case GGML_TYPE_Q6_0_R4: break; - case GGML_TYPE_Q8_0_R4: break; + case GGML_TYPE_Q8_0_R8: break; case GGML_TYPE_Q2_K_R4: break; case GGML_TYPE_Q3_K_R4: break; case GGML_TYPE_Q4_K_R4: break; diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index 4199a282..68525906 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -1606,28 +1606,28 @@ static const ggml_type_traits_t type_traits[GGML_TYPE_COUNT] = { .nrows = 1, .row_meta_size = 0, }, - [GGML_TYPE_IQ4_XS_R4] = { - .type_name = "iq4_xs_r4", + [GGML_TYPE_IQ4_XS_R8] = { + .type_name = "iq4_xs_r8", .blck_size = QK_K, .type_size = sizeof(block_iq4_xs), .is_quantized = true, - .to_float = (ggml_to_float_t) dequantize_row_iq4_xs_r4, - .from_float = quantize_row_iq4_xs_r4, - .from_float_ref = (ggml_from_float_t)quantize_row_iq4_xs_r4_ref, - .vec_dot = vec_dot_iq4_xs_r4_q8_k, + .to_float = (ggml_to_float_t) dequantize_row_iq4_xs_r8, + .from_float = quantize_row_iq4_xs_r8, + .from_float_ref = (ggml_from_float_t)quantize_row_iq4_xs_r8_ref, + .vec_dot = vec_dot_iq4_xs_r8_q8_k, .vec_dot_type = GGML_TYPE_Q8_K32, .nrows = 1, .row_meta_size = 0, }, - [GGML_TYPE_Q4_0_R4] = { - .type_name = "q4_0_r4", + [GGML_TYPE_Q4_0_R8] = { + .type_name = "q4_0_r8", .blck_size = QK4_NL, .type_size = sizeof(block_iq4_nl), .is_quantized = true, - .to_float = (ggml_to_float_t) dequantize_row_q4_0_r4, - .from_float = quantize_row_q4_0_r4, - .from_float_ref = (ggml_from_float_t)quantize_row_q4_0_r4_ref, - .vec_dot = vec_dot_q4_0_r4_q8_0, + .to_float = (ggml_to_float_t) dequantize_row_q4_0_r8, + .from_float = quantize_row_q4_0_r8, + .from_float_ref = (ggml_from_float_t)quantize_row_q4_0_r8_ref, + .vec_dot = vec_dot_q4_0_r8_q8_0, #if GGML_USE_IQK_MULMAT #if defined __AVX2__ .vec_dot_type = GGML_TYPE_Q8_1_X4, @@ -1640,15 +1640,15 @@ static const ggml_type_traits_t type_traits[GGML_TYPE_COUNT] = { .nrows = 1, .row_meta_size = 0, }, - [GGML_TYPE_Q8_0_R4] = { - .type_name = "q8_0_r4", + [GGML_TYPE_Q8_0_R8] = { + .type_name = "q8_0_r8", .blck_size = QK8_0, .type_size = sizeof(block_q8_0), .is_quantized = true, - .to_float = (ggml_to_float_t) dequantize_row_q8_0_r4, - .from_float = quantize_row_q8_0_r4, - .from_float_ref = (ggml_from_float_t)quantize_row_q8_0_r4_ref, - .vec_dot = vec_dot_q8_0_r4_q8_0, + .to_float = (ggml_to_float_t) dequantize_row_q8_0_r8, + .from_float = quantize_row_q8_0_r8, + .from_float_ref = (ggml_from_float_t)quantize_row_q8_0_r8_ref, + .vec_dot = vec_dot_q8_0_r8_q8_0, #if GGML_USE_IQK_MULMAT #if defined __AVX2__ .vec_dot_type = GGML_TYPE_Q8_1_X4, @@ -4390,11 +4390,11 @@ enum ggml_type ggml_ftype_to_ggml_type(enum ggml_ftype ftype) { case GGML_FTYPE_MOSTLY_IQ2_BN_R4: wtype = GGML_TYPE_IQ2_BN_R4;break; case GGML_FTYPE_MOSTLY_IQ4_NL: wtype = GGML_TYPE_IQ4_NL; break; case GGML_FTYPE_MOSTLY_IQ4_NL_R4: wtype = GGML_TYPE_IQ4_NL_R4;break; - case GGML_FTYPE_MOSTLY_IQ4_XS_R4: wtype = GGML_TYPE_IQ4_XS_R4;break; - case GGML_FTYPE_MOSTLY_Q4_0_R4: wtype = GGML_TYPE_Q4_0_R4; break; + case GGML_FTYPE_MOSTLY_IQ4_XS_R8: wtype = GGML_TYPE_IQ4_XS_R8;break; + case GGML_FTYPE_MOSTLY_Q4_0_R8: wtype = GGML_TYPE_Q4_0_R8; break; case GGML_FTYPE_MOSTLY_Q5_0_R4: wtype = GGML_TYPE_Q5_0_R4; break; case GGML_FTYPE_MOSTLY_Q6_0_R4: wtype = GGML_TYPE_Q6_0_R4; break; - case GGML_FTYPE_MOSTLY_Q8_0_R4: wtype = GGML_TYPE_Q8_0_R4; break; + case GGML_FTYPE_MOSTLY_Q8_0_R8: wtype = GGML_TYPE_Q8_0_R8; break; case GGML_FTYPE_MOSTLY_IQ4_XS: wtype = GGML_TYPE_IQ4_XS; break; case GGML_FTYPE_MOSTLY_IQ4_KS: wtype = GGML_TYPE_IQ4_KS; break; case GGML_FTYPE_MOSTLY_IQ4_KS_R4: wtype = GGML_TYPE_IQ4_KS_R4;break; @@ -10938,12 +10938,12 @@ static void ggml_compute_forward_add( case GGML_TYPE_IQ2_BN_R4: case GGML_TYPE_IQ4_NL: case GGML_TYPE_IQ4_NL_R4: - case GGML_TYPE_IQ4_XS_R4: - case GGML_TYPE_Q4_0_R4: + case GGML_TYPE_IQ4_XS_R8: + case GGML_TYPE_Q4_0_R8: case GGML_TYPE_Q5_0_R4: case GGML_TYPE_Q6_0_R4: case GGML_TYPE_I2_S: - case GGML_TYPE_Q8_0_R4: + case GGML_TYPE_Q8_0_R8: case GGML_TYPE_IQ4_XS: case GGML_TYPE_IQ4_KS: case GGML_TYPE_IQ4_KS_R4: @@ -11408,12 +11408,12 @@ static void ggml_compute_forward_add1( case GGML_TYPE_IQ2_BN_R4: case GGML_TYPE_IQ4_NL: case GGML_TYPE_IQ4_NL_R4: - case GGML_TYPE_IQ4_XS_R4: - case GGML_TYPE_Q4_0_R4: + case GGML_TYPE_IQ4_XS_R8: + case GGML_TYPE_Q4_0_R8: case GGML_TYPE_Q5_0_R4: case GGML_TYPE_Q6_0_R4: case GGML_TYPE_I2_S: - case GGML_TYPE_Q8_0_R4: + case GGML_TYPE_Q8_0_R8: case GGML_TYPE_IQ4_XS: case GGML_TYPE_IQ4_KS: case GGML_TYPE_IQ4_KS_R4: @@ -11575,12 +11575,12 @@ static void ggml_compute_forward_acc( case GGML_TYPE_IQ2_BN_R4: case GGML_TYPE_IQ4_NL: case GGML_TYPE_IQ4_NL_R4: - case GGML_TYPE_IQ4_XS_R4: - case GGML_TYPE_Q4_0_R4: + case GGML_TYPE_IQ4_XS_R8: + case GGML_TYPE_Q4_0_R8: case GGML_TYPE_Q5_0_R4: case GGML_TYPE_Q6_0_R4: case GGML_TYPE_I2_S: - case GGML_TYPE_Q8_0_R4: + case GGML_TYPE_Q8_0_R8: case GGML_TYPE_IQ4_XS: case GGML_TYPE_IQ4_KS: case GGML_TYPE_IQ4_KS_R4: @@ -14815,12 +14815,12 @@ static void ggml_compute_forward_out_prod( case GGML_TYPE_IQ2_BN_R4: case GGML_TYPE_IQ4_NL: case GGML_TYPE_IQ4_NL_R4: - case GGML_TYPE_IQ4_XS_R4: - case GGML_TYPE_Q4_0_R4: + case GGML_TYPE_IQ4_XS_R8: + case GGML_TYPE_Q4_0_R8: case GGML_TYPE_Q5_0_R4: case GGML_TYPE_Q6_0_R4: case GGML_TYPE_I2_S: - case GGML_TYPE_Q8_0_R4: + case GGML_TYPE_Q8_0_R8: case GGML_TYPE_IQ4_XS: case GGML_TYPE_IQ4_KS: case GGML_TYPE_IQ4_KS_R4: @@ -15222,12 +15222,12 @@ static void ggml_compute_forward_set( case GGML_TYPE_IQ2_BN_R4: case GGML_TYPE_IQ4_NL: case GGML_TYPE_IQ4_NL_R4: - case GGML_TYPE_IQ4_XS_R4: - case GGML_TYPE_Q4_0_R4: + case GGML_TYPE_IQ4_XS_R8: + case GGML_TYPE_Q4_0_R8: case GGML_TYPE_Q5_0_R4: case GGML_TYPE_Q6_0_R4: case GGML_TYPE_I2_S: - case GGML_TYPE_Q8_0_R4: + case GGML_TYPE_Q8_0_R8: case GGML_TYPE_IQ4_XS: case GGML_TYPE_IQ4_KS: case GGML_TYPE_IQ4_KS_R4: @@ -15523,12 +15523,12 @@ static void ggml_compute_forward_get_rows( case GGML_TYPE_IQ2_BN_R4: case GGML_TYPE_IQ4_NL: case GGML_TYPE_IQ4_NL_R4: - case GGML_TYPE_IQ4_XS_R4: - case GGML_TYPE_Q4_0_R4: + case GGML_TYPE_IQ4_XS_R8: + case GGML_TYPE_Q4_0_R8: case GGML_TYPE_Q5_0_R4: case GGML_TYPE_Q6_0_R4: case GGML_TYPE_I2_S: - case GGML_TYPE_Q8_0_R4: + case GGML_TYPE_Q8_0_R8: case GGML_TYPE_IQ4_XS: case GGML_TYPE_IQ4_KS: case GGML_TYPE_IQ4_KS_R4: @@ -16153,12 +16153,12 @@ static void ggml_compute_forward_clamp( case GGML_TYPE_IQ2_BN_R4: case GGML_TYPE_IQ4_NL: case GGML_TYPE_IQ4_NL_R4: - case GGML_TYPE_IQ4_XS_R4: - case GGML_TYPE_Q4_0_R4: + case GGML_TYPE_IQ4_XS_R8: + case GGML_TYPE_Q4_0_R8: case GGML_TYPE_Q5_0_R4: case GGML_TYPE_Q6_0_R4: case GGML_TYPE_I2_S: - case GGML_TYPE_Q8_0_R4: + case GGML_TYPE_Q8_0_R8: case GGML_TYPE_IQ4_XS: case GGML_TYPE_IQ4_KS: case GGML_TYPE_IQ4_KS_R4: @@ -23028,11 +23028,11 @@ size_t ggml_quantize_chunk( case GGML_TYPE_IQ2_BN_R4:result = quantize_iq2_bn_r4(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_IQ4_NL: result = quantize_iq4_nl (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_IQ4_NL_R4: result = quantize_iq4_nl_r4(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; - case GGML_TYPE_IQ4_XS_R4: result = quantize_iq4_xs_r4(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; - case GGML_TYPE_Q4_0_R4: result = quantize_q4_0_r4(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_IQ4_XS_R8: result = quantize_iq4_xs_r8(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_Q4_0_R8: result = quantize_q4_0_r8(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q5_0_R4: result = quantize_q5_0_r4(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q6_0_R4: result = quantize_q6_0_r4(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; - case GGML_TYPE_Q8_0_R4: result = quantize_q8_0_r4(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_Q8_0_R8: result = quantize_q8_0_r8(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_IQ4_XS: result = quantize_iq4_xs (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_IQ4_KS: result = quantize_iq4_ks (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_IQ4_KS_R4:result = quantize_iq4_ks_r4(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; diff --git a/ggml/src/iqk/iqk_mul_mat.cpp b/ggml/src/iqk/iqk_mul_mat.cpp index 57024602..c561ca2b 100644 --- a/ggml/src/iqk/iqk_mul_mat.cpp +++ b/ggml/src/iqk/iqk_mul_mat.cpp @@ -266,12 +266,12 @@ struct MulMat { case GGML_TYPE_Q5_0_R4: case GGML_TYPE_Q6_0_R4: case GGML_TYPE_IQ2_BN_R4: - case GGML_TYPE_IQ4_XS_R4: + case GGML_TYPE_IQ4_XS_R8: case GGML_TYPE_Q4_K_R4: case GGML_TYPE_Q5_K_R4: case GGML_TYPE_Q8_K_R8: return 8; - case GGML_TYPE_Q4_0_R4: - case GGML_TYPE_Q8_0_R4: + case GGML_TYPE_Q4_0_R8: + case GGML_TYPE_Q8_0_R8: case GGML_TYPE_BF16_R16: return 16; default: return 1; } @@ -298,9 +298,9 @@ struct MulMat { case GGML_TYPE_IQ1_S_R4: case GGML_TYPE_IQ1_M_R4: case GGML_TYPE_IQ2_BN_R4: return 4; - case GGML_TYPE_IQ4_XS_R4: - case GGML_TYPE_Q4_0_R4: - case GGML_TYPE_Q8_0_R4: + case GGML_TYPE_IQ4_XS_R8: + case GGML_TYPE_Q4_0_R8: + case GGML_TYPE_Q8_0_R8: case GGML_TYPE_Q8_K_R8: return 8; case GGML_TYPE_BF16_R16: return 16; default: return 1; @@ -3435,7 +3435,7 @@ inline __m256i accum_q4_0_quants(const __m256i * v, const int8_t * qs) { } template -static void mul_mat_q4_0_r4_q8_1_avx2(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { +static void mul_mat_q4_0_r8_q8_1_avx2(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { GGML_ASSERT(nrc_x%8 == 0); Q8 q8(info); auto m4 = _mm256_set1_epi8(0xf); @@ -3709,9 +3709,9 @@ static void mul_mat_iq1_m_r4_q8_0(int n, const void * vx, size_t bx, const DataI #ifdef HAVE_FANCY_SIMD template -static void mul_mat_q4_0_r4_q8_1(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { +static void mul_mat_q4_0_r8_q8_1(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { if constexpr (nrc_y == 1) { - mul_mat_q4_0_r4_q8_1_avx2<1>(n, vx, bx, info, nrc_x); + mul_mat_q4_0_r8_q8_1_avx2<1>(n, vx, bx, info, nrc_x); return; } GGML_ASSERT(nrc_x%16 == 0); @@ -3787,8 +3787,8 @@ static void mul_mat_q4_0_r4_q8_1(int n, const void * vx, size_t bx, const DataIn } #else template -static void mul_mat_q4_0_r4_q8_1(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { - mul_mat_q4_0_r4_q8_1_avx2(n, vx, bx, info, nrc_x); +static void mul_mat_q4_0_r8_q8_1(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { + mul_mat_q4_0_r8_q8_1_avx2(n, vx, bx, info, nrc_x); } #endif @@ -4177,7 +4177,7 @@ inline __m256i q8_0_r8_dot_product(const uint8_t * x, const int8_t * y, __m256i return qx_r8_q8_dot_product(qx, y); } template -static void mul_mat_q8_0_r4_q8_1(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { +static void mul_mat_q8_0_r8_q8_1(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { GGML_ASSERT(nrc_x%16 == 0); Q8 q8(info); int nb = n / QK8_0; @@ -4263,7 +4263,7 @@ static void mul_mat_q8_0_r4_q8_1(int n, const void * vx, size_t bx, const DataIn } #else template -static void mul_mat_q8_0_r4_q8_1(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { +static void mul_mat_q8_0_r8_q8_1(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { GGML_ASSERT(nrc_x%8 == 0); Q8 q8(info); auto m1 = _mm256_set1_epi16(1); @@ -4345,7 +4345,7 @@ static void mul_mat_q8_0_r4_q8_1(int n, const void * vx, size_t bx, const DataIn #endif template -static void mul_mat_iq4_xs_r4_q8_k_avx2(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { +static void mul_mat_iq4_xs_r8_q8_k_avx2(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { GGML_ASSERT(nrc_x%8 == 0); Q8 q8(info); auto m4 = _mm256_set1_epi8(0xf); @@ -4364,7 +4364,7 @@ static void mul_mat_iq4_xs_r4_q8_k_avx2(int n, const void * vx, size_t bx, const __m256 acc[nrc_y] = {}; __m256i qx[4]; for (int ix = 0; ix < nrc_x; ix += 8) { - const block_iq4_xs_r4 * iq4 = (const block_iq4_xs_r4 *)((const char *)vx + (ix+0)*bx); + const block_iq4_xs_r8 * iq4 = (const block_iq4_xs_r8 *)((const char *)vx + (ix+0)*bx); for (int ibl = 0; ibl < nbl; ++ibl) { // Block of 256 auto d4 = _mm256_cvtph_ps(_mm_loadu_si128((const __m128i *)iq4[ibl].d)); auto slbits = _mm256_loadu_si256((const __m256i *)iq4[ibl].scales_l); @@ -4465,11 +4465,11 @@ static void mul_mat_iq4_xs_r4_q8_k_avx2(int n, const void * vx, size_t bx, const #ifdef HAVE_FANCY_SIMD template -static void mul_mat_iq4_xs_r4_q8_k(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { - mul_mat_iq4_xs_r4_q8_k_avx2(n, vx, bx, info, nrc_x); +static void mul_mat_iq4_xs_r8_q8_k(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { + mul_mat_iq4_xs_r8_q8_k_avx2(n, vx, bx, info, nrc_x); return; if constexpr (nrc_y == 1){ - mul_mat_iq4_xs_r4_q8_k_avx2<1>(n, vx, bx, info, nrc_x); + mul_mat_iq4_xs_r8_q8_k_avx2<1>(n, vx, bx, info, nrc_x); } else { GGML_ASSERT(nrc_x%8 == 0); Q8 q8(info); @@ -4482,8 +4482,8 @@ static void mul_mat_iq4_xs_r4_q8_k(int n, const void * vx, size_t bx, const Data __m512i isum[nrc_y] = {}; __m512i qx[4]; for (int ix = 0; ix < nrc_x; ix += 8) { - const block_iq4_xs_r4 * iq4l = (const block_iq4_xs_r4 *)((const char *)vx + (ix+0)*bx); - const block_iq4_xs_r4 * iq4h = (const block_iq4_xs_r4 *)((const char *)vx + (ix+4)*bx); + const block_iq4_xs_r8 * iq4l = (const block_iq4_xs_r8 *)((const char *)vx + (ix+0)*bx); + const block_iq4_xs_r8 * iq4h = (const block_iq4_xs_r8 *)((const char *)vx + (ix+4)*bx); for (int ibl = 0; ibl < nbl; ++ibl) { // Block of 256 auto dl = _mm_cvtph_ps(_mm_loadl_epi64((const __m128i *)iq4l[ibl].d)); auto dh = _mm_cvtph_ps(_mm_loadl_epi64((const __m128i *)iq4h[ibl].d)); @@ -4544,8 +4544,8 @@ static void mul_mat_iq4_xs_r4_q8_k(int n, const void * vx, size_t bx, const Data } #else template -static void mul_mat_iq4_xs_r4_q8_k(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { - mul_mat_iq4_xs_r4_q8_k_avx2(n, vx, bx, info, nrc_x); +static void mul_mat_iq4_xs_r8_q8_k(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { + mul_mat_iq4_xs_r8_q8_k_avx2(n, vx, bx, info, nrc_x); } #endif @@ -8889,16 +8889,16 @@ bool MulMat::prepare(int typeA, int typeB, int ne00, MulMat& mm, int Ny) { mm.funcs[7] = mul_mat_iq4_nl_r4_q8_1<8>; expected_typeB = GGML_TYPE_Q8_1_X4; break; - case GGML_TYPE_IQ4_XS_R4: + case GGML_TYPE_IQ4_XS_R8: assert (ne00 % QK_K == 0); - mm.funcs[0] = mul_mat_iq4_xs_r4_q8_k<1>; - mm.funcs[1] = mul_mat_iq4_xs_r4_q8_k<2>; - mm.funcs[2] = mul_mat_iq4_xs_r4_q8_k<3>; - mm.funcs[3] = mul_mat_iq4_xs_r4_q8_k<4>; - mm.funcs[4] = mul_mat_iq4_xs_r4_q8_k<5>; - mm.funcs[5] = mul_mat_iq4_xs_r4_q8_k<6>; - mm.funcs[6] = mul_mat_iq4_xs_r4_q8_k<7>; - mm.funcs[7] = mul_mat_iq4_xs_r4_q8_k<8>; + mm.funcs[0] = mul_mat_iq4_xs_r8_q8_k<1>; + mm.funcs[1] = mul_mat_iq4_xs_r8_q8_k<2>; + mm.funcs[2] = mul_mat_iq4_xs_r8_q8_k<3>; + mm.funcs[3] = mul_mat_iq4_xs_r8_q8_k<4>; + mm.funcs[4] = mul_mat_iq4_xs_r8_q8_k<5>; + mm.funcs[5] = mul_mat_iq4_xs_r8_q8_k<6>; + mm.funcs[6] = mul_mat_iq4_xs_r8_q8_k<7>; + mm.funcs[7] = mul_mat_iq4_xs_r8_q8_k<8>; expected_typeB = GGML_TYPE_Q8_K32; break; case GGML_TYPE_IQ4_KS_R4: @@ -9113,18 +9113,18 @@ bool MulMat::prepare(int typeA, int typeB, int ne00, MulMat& mm, int Ny) { #endif expected_typeB = GGML_TYPE_Q8_K; break; - case GGML_TYPE_Q4_0_R4: + case GGML_TYPE_Q4_0_R8: assert (ne00 % QK4_NL == 0); - mm.funcs[0] = mul_mat_q4_0_r4_q8_1<1>; - mm.funcs[1] = mul_mat_q4_0_r4_q8_1<2>; - mm.funcs[2] = mul_mat_q4_0_r4_q8_1<3>; - mm.funcs[3] = mul_mat_q4_0_r4_q8_1<4>; - mm.funcs[4] = mul_mat_q4_0_r4_q8_1<5>; - mm.funcs[5] = mul_mat_q4_0_r4_q8_1<6>; - mm.funcs[6] = mul_mat_q4_0_r4_q8_1<7>; - mm.funcs[7] = mul_mat_q4_0_r4_q8_1<8>; + mm.funcs[0] = mul_mat_q4_0_r8_q8_1<1>; + mm.funcs[1] = mul_mat_q4_0_r8_q8_1<2>; + mm.funcs[2] = mul_mat_q4_0_r8_q8_1<3>; + mm.funcs[3] = mul_mat_q4_0_r8_q8_1<4>; + mm.funcs[4] = mul_mat_q4_0_r8_q8_1<5>; + mm.funcs[5] = mul_mat_q4_0_r8_q8_1<6>; + mm.funcs[6] = mul_mat_q4_0_r8_q8_1<7>; + mm.funcs[7] = mul_mat_q4_0_r8_q8_1<8>; #ifdef HAVE_FANCY_SIMD - mm.func16 = mul_mat_q4_0_r4_q8_1<16>; + mm.func16 = mul_mat_q4_0_r8_q8_1<16>; #endif expected_typeB = GGML_TYPE_Q8_1_X4; break; @@ -9152,16 +9152,16 @@ bool MulMat::prepare(int typeA, int typeB, int ne00, MulMat& mm, int Ny) { mm.funcs[7] = mul_mat_q6_0_r4_q8_1<8>; expected_typeB = GGML_TYPE_Q8_1_X4; break; - case GGML_TYPE_Q8_0_R4: + case GGML_TYPE_Q8_0_R8: assert (ne00 % QK4_NL == 0); - mm.funcs[0] = mul_mat_q8_0_r4_q8_1<1>; - mm.funcs[1] = mul_mat_q8_0_r4_q8_1<2>; - mm.funcs[2] = mul_mat_q8_0_r4_q8_1<3>; - mm.funcs[3] = mul_mat_q8_0_r4_q8_1<4>; - mm.funcs[4] = mul_mat_q8_0_r4_q8_1<5>; - mm.funcs[5] = mul_mat_q8_0_r4_q8_1<6>; - mm.funcs[6] = mul_mat_q8_0_r4_q8_1<7>; - mm.funcs[7] = mul_mat_q8_0_r4_q8_1<8>; + mm.funcs[0] = mul_mat_q8_0_r8_q8_1<1>; + mm.funcs[1] = mul_mat_q8_0_r8_q8_1<2>; + mm.funcs[2] = mul_mat_q8_0_r8_q8_1<3>; + mm.funcs[3] = mul_mat_q8_0_r8_q8_1<4>; + mm.funcs[4] = mul_mat_q8_0_r8_q8_1<5>; + mm.funcs[5] = mul_mat_q8_0_r8_q8_1<6>; + mm.funcs[6] = mul_mat_q8_0_r8_q8_1<7>; + mm.funcs[7] = mul_mat_q8_0_r8_q8_1<8>; expected_typeB = GGML_TYPE_Q8_1_X4; break; case GGML_TYPE_IQ1_S_R4: @@ -11779,7 +11779,7 @@ IQK_ALWAYS_INLINE void prepare_iq4_nl_quants_r8(const int8x16_t& values, const u } template -void mul_mat_iq4_xs_r4_q8_k(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { +void mul_mat_iq4_xs_r8_q8_k(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { GGML_ASSERT(nrc_x%4 == 0); Q8 q8(info); auto m4 = vdupq_n_u8(0xf); @@ -11792,7 +11792,7 @@ void mul_mat_iq4_xs_r4_q8_k(int n, const void * vx, size_t bx, const DataInfo& i int32x4x2_t scales; float32x4_t acc[2*nrc_y] = {}; for (int ix = 0; ix < nrc_x; ix += 8) { - const block_iq4_xs_r4 * iq4 = (const block_iq4_xs_r4 *)((const char *)vx + ix*bx); + const block_iq4_xs_r8 * iq4 = (const block_iq4_xs_r8 *)((const char *)vx + ix*bx); for (int ibl = 0; ibl < nbl; ++ibl) { auto d4_f16 = vld1q_f16((const float16_t *)iq4[ibl].d); auto d4l = vcvt_f32_f16(vget_low_f16 (d4_f16)); @@ -13662,7 +13662,7 @@ inline void qx_0_q8_0_dot(const int8x16_t * qx, const int8_t * qy, int32x4_t& su } template -void mul_mat_q8_0_r4_q8_0(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { +void mul_mat_q8_0_r8_q8_0(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { GGML_ASSERT(nrc_x%8 == 0); Q8 q8(info); int nb = n / QK8_0; @@ -13880,8 +13880,8 @@ bool MulMat::prepare(int typeA, int typeB, int ne00, MulMat& m, int /*Ny*/) { SET_MUL_MAT_FUNCTIONS_T(m, mul_mat_qx_r4_q8_0, IQ4_NL_R4_Dequantizer); expected_Btype = GGML_TYPE_Q8_0_X4; break; - case GGML_TYPE_IQ4_XS_R4: - SET_MUL_MAT_FUNCTIONS(m, mul_mat_iq4_xs_r4_q8_k); + case GGML_TYPE_IQ4_XS_R8: + SET_MUL_MAT_FUNCTIONS(m, mul_mat_iq4_xs_r8_q8_k); expected_Btype = GGML_TYPE_Q8_K32; break; case GGML_TYPE_IQ4_KS_R4: @@ -13964,7 +13964,7 @@ bool MulMat::prepare(int typeA, int typeB, int ne00, MulMat& m, int /*Ny*/) { SET_MUL_MAT_FUNCTIONS(m, mul_mat_iq5_k_r4_q8_k); expected_Btype = GGML_TYPE_Q8_K; break; - case GGML_TYPE_Q4_0_R4: + case GGML_TYPE_Q4_0_R8: SET_MUL_MAT_FUNCTIONS_T(m, mul_mat_qx_r8_q8_0, Q4_0_R8_Dequantizer); expected_Btype = GGML_TYPE_Q8_0_X4; break; @@ -13976,8 +13976,8 @@ bool MulMat::prepare(int typeA, int typeB, int ne00, MulMat& m, int /*Ny*/) { SET_MUL_MAT_FUNCTIONS_T(m, mul_mat_qx_r4_q8_0, Q6_0_R4_Dequantizer); expected_Btype = GGML_TYPE_Q8_0_X4; break; - case GGML_TYPE_Q8_0_R4: - SET_MUL_MAT_FUNCTIONS(m, mul_mat_q8_0_r4_q8_0); + case GGML_TYPE_Q8_0_R8: + SET_MUL_MAT_FUNCTIONS(m, mul_mat_q8_0_r8_q8_0); expected_Btype = GGML_TYPE_Q8_0_X4; break; default: @@ -15260,9 +15260,9 @@ struct FlashQKfp32 { } else if constexpr (std::is_same_v>) { #ifdef __aarch64__ - MAKE_FUNCS_ONLY_NRC(mul_mat_q8_0_r4_q8_0, nq); + MAKE_FUNCS_ONLY_NRC(mul_mat_q8_0_r8_q8_0, nq); #else - MAKE_FUNCS_ONLY_NRC(mul_mat_q8_0_r4_q8_1, nq); + MAKE_FUNCS_ONLY_NRC(mul_mat_q8_0_r8_q8_1, nq); #endif } else if constexpr (std::is_same_v>) { diff --git a/ggml/src/iqk/iqk_quantize.cpp b/ggml/src/iqk/iqk_quantize.cpp index e741a8ea..9ce5731d 100644 --- a/ggml/src/iqk/iqk_quantize.cpp +++ b/ggml/src/iqk/iqk_quantize.cpp @@ -3622,16 +3622,16 @@ void vec_dot_iq4_nl_r4_q8_0(int n, float * s, size_t bs, const void * vx, size_t } // -// ========================================= q4_0_r4 +// ========================================= q4_0_r8 // -void quantize_row_q4_0_r4_ref(const float * x, block_iq4_nl_r8 * y, int64_t k) { +void quantize_row_q4_0_r8_ref(const float * x, block_iq4_nl_r8 * y, int64_t k) { // we assume we are called with 8 rows - quantize_q4_0_r4(x, (void *)y, 8, k/8, nullptr); + quantize_q4_0_r8(x, (void *)y, 8, k/8, nullptr); } -void quantize_row_q4_0_r4(const float * x, void * y, int64_t k) { +void quantize_row_q4_0_r8(const float * x, void * y, int64_t k) { // we assume we are called with 8 rows - quantize_q4_0_r4(x, y, 8, k/8, nullptr); + quantize_q4_0_r8(x, y, 8, k/8, nullptr); } static void repack_q4_0(int nrows, int n_per_row, const block_q4_0 * x, block_iq4_nl_r8 * y, [[maybe_unused]] bool online) { @@ -3664,7 +3664,7 @@ static void repack_q4_0(int nrows, int n_per_row, const block_q4_0 * x, block_iq } } #ifdef __ARM_NEON -static void modify_q4_0_r4(int64_t k, char * cy) { +static void modify_q4_0_r8(int64_t k, char * cy) { auto y = (block_iq4_nl_r8 *)cy; int nb = k/(32*8); for (int ib = 0; ib < nb; ++ib) { @@ -3680,7 +3680,7 @@ static void modify_q4_0_r4(int64_t k, char * cy) { } #endif -size_t quantize_q4_0_r4(const float * src, void * dst, int64_t nrows, int64_t n_per_row, const float * imatrix) { +size_t quantize_q4_0_r8(const float * src, void * dst, int64_t nrows, int64_t n_per_row, const float * imatrix) { GGML_ASSERT(nrows%8 == 0); auto row_size_nl = ggml_row_size(GGML_TYPE_IQ4_NL, n_per_row); std::vector qtmp(8*row_size_nl); @@ -3694,7 +3694,7 @@ size_t quantize_q4_0_r4(const float * src, void * dst, int64_t nrows, int64_t n_ return nrows*row_size_nl; } -void dequantize_row_q4_0_r4(const block_iq4_nl_r8 * x, float * y, int64_t k) { +void dequantize_row_q4_0_r8(const block_iq4_nl_r8 * x, float * y, int64_t k) { // we assume we are called with 8 rows int n_per_row = k/8; int nb = n_per_row/QK4_0; @@ -3713,9 +3713,9 @@ void dequantize_row_q4_0_r4(const block_iq4_nl_r8 * x, float * y, int64_t k) { } } -void vec_dot_q4_0_r4_q8_0(int n, float * s, size_t bs, const void * vx, size_t bx, const void * vy, size_t by, int nrc) { +void vec_dot_q4_0_r8_q8_0(int n, float * s, size_t bs, const void * vx, size_t bx, const void * vy, size_t by, int nrc) { #if GGML_USE_IQK_MULMAT - if (iqk_mul_mat(1, 1, n, GGML_TYPE_Q4_0_R4, vx, 0, GGML_TYPE_Q8_0, vy, 0, s, 0, 0, 1)) { + if (iqk_mul_mat(1, 1, n, GGML_TYPE_Q4_0_R8, vx, 0, GGML_TYPE_Q8_0, vy, 0, s, 0, 0, 1)) { return; } #endif @@ -3728,16 +3728,16 @@ void vec_dot_q4_0_r4_q8_0(int n, float * s, size_t bs, const void * vx, size_t b // -// ========================================= q8_0_r4 +// ========================================= q8_0_r8 // -void quantize_row_q8_0_r4_ref(const float * x, block_q8_0_r8 * y, int64_t k) { +void quantize_row_q8_0_r8_ref(const float * x, block_q8_0_r8 * y, int64_t k) { // we assume we are called with 4 rows - quantize_q8_0_r4(x, (void *)y, 8, k/8, nullptr); + quantize_q8_0_r8(x, (void *)y, 8, k/8, nullptr); } -void quantize_row_q8_0_r4(const float * x, void * y, int64_t k) { +void quantize_row_q8_0_r8(const float * x, void * y, int64_t k) { // we assume we are called with 4 rows - quantize_q8_0_r4(x, y, 8, k/8, nullptr); + quantize_q8_0_r8(x, y, 8, k/8, nullptr); } static void repack_q8_0(int nrows, int n_per_row, const block_q8_0 * x, block_q8_0_r8 * y, [[maybe_unused]] bool online) { @@ -3770,7 +3770,7 @@ static void repack_q8_0(int nrows, int n_per_row, const block_q8_0 * x, block_q8 } #ifdef HAVE_FANCY_SIMD -static void modify_q8_0_r4(int64_t k, char * cy) { +static void modify_q8_0_r8(int64_t k, char * cy) { auto y = (block_iq4_nl_r8 *)cy; int nb = k/(32*8); for (int ib = 0; ib < nb; ++ib) { @@ -3782,7 +3782,7 @@ static void modify_q8_0_r4(int64_t k, char * cy) { } #endif -size_t quantize_q8_0_r4(const float * src, void * dst, int64_t nrows, int64_t n_per_row, const float * imatrix) { +size_t quantize_q8_0_r8(const float * src, void * dst, int64_t nrows, int64_t n_per_row, const float * imatrix) { GGML_ASSERT(nrows%8 == 0); auto row_size_0 = ggml_row_size(GGML_TYPE_Q8_0, n_per_row); std::vector qtmp(8*row_size_0); @@ -3796,7 +3796,7 @@ size_t quantize_q8_0_r4(const float * src, void * dst, int64_t nrows, int64_t n_ return nrows*row_size_0; } -void dequantize_row_q8_0_r4(const block_q8_0_r8 * x, float * y, int64_t k) { +void dequantize_row_q8_0_r8(const block_q8_0_r8 * x, float * y, int64_t k) { // we assume we are called with 4 rows int n_per_row = k/8; int nb = n_per_row/QK8_0; @@ -3813,9 +3813,9 @@ void dequantize_row_q8_0_r4(const block_q8_0_r8 * x, float * y, int64_t k) { } } -void vec_dot_q8_0_r4_q8_0(int n, float * s, size_t bs, const void * vx, size_t bx, const void * vy, size_t by, int nrc) { +void vec_dot_q8_0_r8_q8_0(int n, float * s, size_t bs, const void * vx, size_t bx, const void * vy, size_t by, int nrc) { #if GGML_USE_IQK_MULMAT - if (iqk_mul_mat(1, 1, n, GGML_TYPE_Q8_0_R4, vx, 0, GGML_TYPE_Q8_0, vy, 0, s, 0, 0, 1)) { + if (iqk_mul_mat(1, 1, n, GGML_TYPE_Q8_0_R8, vx, 0, GGML_TYPE_Q8_0, vy, 0, s, 0, 0, 1)) { return; } #endif @@ -4025,18 +4025,18 @@ void vec_dot_q6_0_r4_q8_0(int n, float * s, size_t bs, const void * vx, size_t b } // -// ========================================= iq4_xs_r4 +// ========================================= iq4_xs_r8 // -void quantize_row_iq4_xs_r4_ref(const float * x, block_iq4_xs_r4 * y, int64_t k) { - quantize_iq4_xs_r4(x, (void *)y, 8, k/8, nullptr); +void quantize_row_iq4_xs_r8_ref(const float * x, block_iq4_xs_r8 * y, int64_t k) { + quantize_iq4_xs_r8(x, (void *)y, 8, k/8, nullptr); } -void quantize_row_iq4_xs_r4(const float * x, void * y, int64_t k) { - quantize_iq4_xs_r4(x, y, 8, k/8, nullptr); +void quantize_row_iq4_xs_r8(const float * x, void * y, int64_t k) { + quantize_iq4_xs_r8(x, y, 8, k/8, nullptr); } -static void repack_iq4_xs(int nrows, int n_per_row, const block_iq4_xs * x, block_iq4_xs_r4 * y, [[maybe_unused]] bool online) { +static void repack_iq4_xs(int nrows, int n_per_row, const block_iq4_xs * x, block_iq4_xs_r8 * y, [[maybe_unused]] bool online) { GGML_ASSERT(nrows%8 == 0); GGML_ASSERT(n_per_row%QK_K == 0); int nblock = n_per_row/QK_K; @@ -4068,7 +4068,7 @@ static void repack_iq4_xs(int nrows, int n_per_row, const block_iq4_xs * x, bloc } } -size_t quantize_iq4_xs_r4(const float * src, void * dst, int64_t nrows, int64_t n_per_row, const float * imatrix) { +size_t quantize_iq4_xs_r8(const float * src, void * dst, int64_t nrows, int64_t n_per_row, const float * imatrix) { GGML_ASSERT(nrows%8 == 0); GGML_ASSERT(n_per_row%QK_K == 0); char * qcur = (char *)dst; @@ -4076,14 +4076,14 @@ size_t quantize_iq4_xs_r4(const float * src, void * dst, int64_t nrows, int64_t std::vector qtmp(8*row_size); for (int row = 0; row < nrows; row += 8) { quantize_iq4_xs(src, (void *)qtmp.data(), 8, n_per_row, imatrix); - repack_iq4_xs(8, n_per_row, (const block_iq4_xs *)qtmp.data(), (block_iq4_xs_r4 *)qcur, false); + repack_iq4_xs(8, n_per_row, (const block_iq4_xs *)qtmp.data(), (block_iq4_xs_r8 *)qcur, false); qcur += 8*row_size; src += 8*n_per_row; } return nrows*row_size; } -void dequantize_row_iq4_xs_r4(const block_iq4_xs_r4 * x, float * y, int64_t k) { +void dequantize_row_iq4_xs_r8(const block_iq4_xs_r8 * x, float * y, int64_t k) { auto n_per_row = k/8; float * y8[8]; for (int k = 0; k < 8; ++k) y8[k] = y + n_per_row*k; @@ -4103,9 +4103,9 @@ void dequantize_row_iq4_xs_r4(const block_iq4_xs_r4 * x, float * y, int64_t k) { } } -void vec_dot_iq4_xs_r4_q8_k(int n, float * s, size_t bs, const void * vx, size_t bx, const void * vy, size_t by, int nrc) { +void vec_dot_iq4_xs_r8_q8_k(int n, float * s, size_t bs, const void * vx, size_t bx, const void * vy, size_t by, int nrc) { #if GGML_USE_IQK_MULMAT - if (iqk_mul_mat(1, 1, n, GGML_TYPE_IQ4_XS_R4, vx, 0, GGML_TYPE_Q8_K, vy, 0, s, 0, 0, 1)) { + if (iqk_mul_mat(1, 1, n, GGML_TYPE_IQ4_XS_R8, vx, 0, GGML_TYPE_Q8_K, vy, 0, s, 0, 0, 1)) { return; } #endif @@ -6329,10 +6329,10 @@ struct Modify { bool iqk_modify_tensor(struct ggml_tensor * tensor) { static const std::unordered_map k_mod_map = { #ifdef __ARM_NEON - { GGML_TYPE_Q4_0_R4, {modify_q4_0_r4, 8} }, + { GGML_TYPE_Q4_0_R8, {modify_q4_0_r8, 8} }, #endif #ifdef HAVE_FANCY_SIMD - { GGML_TYPE_Q8_0_R4, {modify_q8_0_r4, 8} }, + { GGML_TYPE_Q8_0_R8, {modify_q8_0_r8, 8} }, { GGML_TYPE_Q8_K_R8, {modify_q8_k_r8, 8} }, #endif }; @@ -6373,7 +6373,7 @@ void iqk_repack_tensor(struct ggml_tensor * tensor) { { GGML_TYPE_IQ3_K, { GGML_TYPE_IQ3_K_R4, 4, (Repack::repack_func)repack_iq3_k} }, { GGML_TYPE_IQ4_K, { GGML_TYPE_IQ4_K_R4, 4, (Repack::repack_func)repack_iq4_k} }, { GGML_TYPE_IQ5_K, { GGML_TYPE_IQ5_K_R4, 4, (Repack::repack_func)repack_iq5_k} }, - { GGML_TYPE_IQ4_XS, { GGML_TYPE_IQ4_XS_R4, 8, (Repack::repack_func)repack_iq4_xs} }, + { GGML_TYPE_IQ4_XS, { GGML_TYPE_IQ4_XS_R8, 8, (Repack::repack_func)repack_iq4_xs} }, { GGML_TYPE_IQ4_KS, { GGML_TYPE_IQ4_KS_R4, 4, (Repack::repack_func)repack_iq4_ks} }, { GGML_TYPE_IQ4_NL, { GGML_TYPE_IQ4_NL_R4, 4, (Repack::repack_func)repack_iq4_nl} }, { GGML_TYPE_IQ2_BN, { GGML_TYPE_IQ2_BN_R4, 4, (Repack::repack_func)repack_iq2_bn} }, @@ -6387,10 +6387,10 @@ void iqk_repack_tensor(struct ggml_tensor * tensor) { { GGML_TYPE_Q4_K, { GGML_TYPE_Q4_K_R4, 4, (Repack::repack_func)repack_q4_k} }, { GGML_TYPE_Q5_K, { GGML_TYPE_Q5_K_R4, 4, (Repack::repack_func)repack_q5_k} }, { GGML_TYPE_Q6_K, { GGML_TYPE_Q6_K_R4, 4, (Repack::repack_func)repack_q6_k} }, - { GGML_TYPE_Q4_0, { GGML_TYPE_Q4_0_R4, 8, (Repack::repack_func)repack_q4_0} }, + { GGML_TYPE_Q4_0, { GGML_TYPE_Q4_0_R8, 8, (Repack::repack_func)repack_q4_0} }, { GGML_TYPE_Q5_0, { GGML_TYPE_Q5_0_R4, 4, (Repack::repack_func)repack_q5_0} }, { GGML_TYPE_Q6_0, { GGML_TYPE_Q6_0_R4, 4, (Repack::repack_func)repack_q6_0} }, - { GGML_TYPE_Q8_0, { GGML_TYPE_Q8_0_R4, 8, (Repack::repack_func)repack_q8_0} }, + { GGML_TYPE_Q8_0, { GGML_TYPE_Q8_0_R8, 8, (Repack::repack_func)repack_q8_0} }, { GGML_TYPE_Q8_K, { GGML_TYPE_Q8_K_R8, 8, (Repack::repack_func)repack_q8_k} }, #ifdef __AVX512BF16__ { GGML_TYPE_BF16, { GGML_TYPE_BF16_R16, 16, (Repack::repack_func)repack_bf16}}, diff --git a/ggml/src/iqk/iqk_quantize.h b/ggml/src/iqk/iqk_quantize.h index 0dbb88bd..ff553ae7 100644 --- a/ggml/src/iqk/iqk_quantize.h +++ b/ggml/src/iqk/iqk_quantize.h @@ -67,17 +67,17 @@ size_t quantize_iq4_nl_r4(const float * GGML_RESTRICT src, void * GGML_RESTRICT void dequantize_row_iq4_nl_r4(const block_iq4_nl_r4 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); void vec_dot_iq4_nl_r4_q8_0(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 quantize_row_q4_0_r4_ref(const float * GGML_RESTRICT x, block_iq4_nl_r8 * GGML_RESTRICT y, int64_t k); -void quantize_row_q4_0_r4(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); -size_t quantize_q4_0_r4(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); -void dequantize_row_q4_0_r4(const block_iq4_nl_r8 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); -void vec_dot_q4_0_r4_q8_0(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 quantize_row_q8_0_r4_ref(const float * GGML_RESTRICT x, block_q8_0_r8 * GGML_RESTRICT y, int64_t k); -void quantize_row_q8_0_r4(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); -size_t quantize_q8_0_r4(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); -void dequantize_row_q8_0_r4(const block_q8_0_r8 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); -void vec_dot_q8_0_r4_q8_0(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 quantize_row_q4_0_r8_ref(const float * GGML_RESTRICT x, block_iq4_nl_r8 * GGML_RESTRICT y, int64_t k); +void quantize_row_q4_0_r8(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +size_t quantize_q4_0_r8(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); +void dequantize_row_q4_0_r8(const block_iq4_nl_r8 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +void vec_dot_q4_0_r8_q8_0(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 quantize_row_q8_0_r8_ref(const float * GGML_RESTRICT x, block_q8_0_r8 * GGML_RESTRICT y, int64_t k); +void quantize_row_q8_0_r8(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +size_t quantize_q8_0_r8(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); +void dequantize_row_q8_0_r8(const block_q8_0_r8 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +void vec_dot_q8_0_r8_q8_0(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 quantize_row_q5_0_r4_ref(const float * GGML_RESTRICT x, block_q5_0_r4 * GGML_RESTRICT y, int64_t k); void quantize_row_q5_0_r4(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); @@ -91,11 +91,11 @@ size_t quantize_q6_0_r4(const float * GGML_RESTRICT src, void * GGML_RESTRICT ds void dequantize_row_q6_0_r4(const block_q6_0_r4 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); void vec_dot_q6_0_r4_q8_0(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 quantize_row_iq4_xs_r4_ref(const float * GGML_RESTRICT x, block_iq4_xs_r4 * GGML_RESTRICT y, int64_t k); -void quantize_row_iq4_xs_r4(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); -size_t quantize_iq4_xs_r4(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); -void dequantize_row_iq4_xs_r4(const block_iq4_xs_r4 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); -void vec_dot_iq4_xs_r4_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 quantize_row_iq4_xs_r8_ref(const float * GGML_RESTRICT x, block_iq4_xs_r8 * GGML_RESTRICT y, int64_t k); +void quantize_row_iq4_xs_r8(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +size_t quantize_iq4_xs_r8(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); +void dequantize_row_iq4_xs_r8(const block_iq4_xs_r8 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +void vec_dot_iq4_xs_r8_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 quantize_row_iq2_bn_ref (const float * GGML_RESTRICT x, block_iq2_bn * GGML_RESTRICT y, int64_t k); void quantize_row_iq2_bn (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); -- cgit v1.2.3