diff options
| -rw-r--r-- | examples/quantize/quantize.cpp | 2 | ||||
| -rw-r--r-- | ggml/include/ggml.h | 4 | ||||
| -rw-r--r-- | ggml/src/ggml-common.h | 4 | ||||
| -rw-r--r-- | ggml/src/ggml-quants.c | 2 | ||||
| -rw-r--r-- | ggml/src/ggml.c | 30 | ||||
| -rw-r--r-- | ggml/src/iqk/iqk_mul_mat.cpp | 50 | ||||
| -rw-r--r-- | ggml/src/iqk/iqk_quantize.cpp | 30 | ||||
| -rw-r--r-- | ggml/src/iqk/iqk_quantize.h | 14 | ||||
| -rw-r--r-- | include/llama.h | 2 | ||||
| -rw-r--r-- | src/llama.cpp | 20 |
10 files changed, 79 insertions, 79 deletions
diff --git a/examples/quantize/quantize.cpp b/examples/quantize/quantize.cpp index 89fb8464..24e5a61b 100644 --- a/examples/quantize/quantize.cpp +++ b/examples/quantize/quantize.cpp @@ -41,7 +41,7 @@ static const std::vector<struct quant_option> QUANT_OPTIONS = { { "Q3_K_M", LLAMA_FTYPE_MOSTLY_Q3_K_M, " 3.07G, +0.2496 ppl @ LLaMA-v1-7B", }, { "Q3_K_L", LLAMA_FTYPE_MOSTLY_Q3_K_L, " 3.35G, +0.1764 ppl @ LLaMA-v1-7B", }, { "IQ4_NL", LLAMA_FTYPE_MOSTLY_IQ4_NL, " 4.50 bpw non-linear quantization", }, - { "IQ4_NL_X4",LLAMA_FTYPE_MOSTLY_IQ4_NL_X4," 4.50 bpw non-linear quantization", }, + { "IQ4_NL_R4",LLAMA_FTYPE_MOSTLY_IQ4_NL_R4," 4.50 bpw non-linear quantization", }, { "IQ4_XS_R4",LLAMA_FTYPE_MOSTLY_IQ4_XS_R4," 4.25 bpw non-linear quantization", }, { "Q4_0_R4", LLAMA_FTYPE_MOSTLY_Q4_0_R4, " 4.50 bpw quantization", }, { "Q5_0_R4", LLAMA_FTYPE_MOSTLY_Q5_0_R4, " 5.50 bpw quantization", }, diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h index 962e9688..e17540e9 100644 --- a/ggml/include/ggml.h +++ b/ggml/include/ggml.h @@ -411,7 +411,7 @@ extern "C" { GGML_TYPE_Q4_0_R4 = 202, GGML_TYPE_Q5_0_R4 = 206, GGML_TYPE_Q8_0_R4 = 208, - GGML_TYPE_IQ4_NL_X4 = 220, // TODO: rename GGML_TYPE_IQ4_NL_X4 to GGML_TYPE_IQ4_NL_R4 + GGML_TYPE_IQ4_NL_R4 = 220, GGML_TYPE_IQ4_XS_R4 = 223, GGML_TYPE_Q6_0_R4 = 233, GGML_TYPE_IQ2_BN_R4 = 335, @@ -477,7 +477,7 @@ extern "C" { GGML_FTYPE_MOSTLY_Q4_0_R4 = 202, // except 1d tensors GGML_FTYPE_MOSTLY_Q8_0_R4 = 207, // except 1d tensors GGML_FTYPE_MOSTLY_Q5_0_R4 = 208, // except 1d tensors - GGML_FTYPE_MOSTLY_IQ4_NL_X4 = 219, // except 1d tensors + GGML_FTYPE_MOSTLY_IQ4_NL_R4 = 219, // except 1d tensors GGML_FTYPE_MOSTLY_IQ4_XS_R4 = 222, // except 1d tensors GGML_FTYPE_MOSTLY_Q6_0_R4 = 227, // except 1d tensors GGML_FTYPE_MOSTLY_IQ2_BN_R4 = 329, // except 1d tensors diff --git a/ggml/src/ggml-common.h b/ggml/src/ggml-common.h index aa41bf55..3f701873 100644 --- a/ggml/src/ggml-common.h +++ b/ggml/src/ggml-common.h @@ -436,8 +436,8 @@ static_assert(sizeof(block_iq4_nl) == sizeof(ggml_half) + QK4_NL/2, "wrong iq4_n typedef struct { ggml_half d[4]; uint8_t qs[2*QK4_NL]; -} block_iq4_nl_x4; -static_assert(sizeof(block_iq4_nl_x4) == 4*sizeof(ggml_half) + 2*QK4_NL, "wrong iq4_nl_x4 block size/padding"); +} block_iq4_nl_r4; +static_assert(sizeof(block_iq4_nl_r4) == 4*sizeof(ggml_half) + 2*QK4_NL, "wrong iq4_nl_x4 block size/padding"); typedef struct { ggml_half d; diff --git a/ggml/src/ggml-quants.c b/ggml/src/ggml-quants.c index a9511f00..7eece2b3 100644 --- a/ggml/src/ggml-quants.c +++ b/ggml/src/ggml-quants.c @@ -15196,7 +15196,7 @@ bool ggml_validate_row_data(enum ggml_type type, const void * data, size_t nbyte case GGML_TYPE_IQ6_K: break; case GGML_TYPE_IQ4_KS: break; case GGML_TYPE_IQ4_KSS: break; - case GGML_TYPE_IQ4_NL_X4: break; + case GGML_TYPE_IQ4_NL_R4: break; case GGML_TYPE_IQ4_XS_R4: break; case GGML_TYPE_Q4_0_R4: break; case GGML_TYPE_Q5_0_R4: break; diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index 12afec05..69bb6d88 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -1266,15 +1266,15 @@ static const ggml_type_traits_t type_traits[GGML_TYPE_COUNT] = { .nrows = 1, .row_meta_size = 0, }, - [GGML_TYPE_IQ4_NL_X4] = { - .type_name = "iq4_nl_x4", + [GGML_TYPE_IQ4_NL_R4] = { + .type_name = "iq4_nl_r4", .blck_size = QK4_NL, .type_size = sizeof(block_iq4_nl), .is_quantized = true, - .to_float = (ggml_to_float_t) dequantize_row_iq4_nl_x4, - .from_float = quantize_row_iq4_nl_x4, - .from_float_ref = (ggml_from_float_t)quantize_row_iq4_nl_x4_ref, - .vec_dot = vec_dot_iq4_nl_x4_q8_0, + .to_float = (ggml_to_float_t) dequantize_row_iq4_nl_r4, + .from_float = quantize_row_iq4_nl_r4, + .from_float_ref = (ggml_from_float_t)quantize_row_iq4_nl_r4_ref, + .vec_dot = vec_dot_iq4_nl_r4_q8_0, #if GGML_USE_IQK_MULMAT && defined __AVX2__ .vec_dot_type = GGML_TYPE_Q8_1, #else @@ -4023,7 +4023,7 @@ enum ggml_type ggml_ftype_to_ggml_type(enum ggml_ftype ftype) { case GGML_FTYPE_MOSTLY_IQ2_BN: wtype = GGML_TYPE_IQ2_BN; break; 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_X4: wtype = GGML_TYPE_IQ4_NL_X4;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_Q5_0_R4: wtype = GGML_TYPE_Q5_0_R4; break; @@ -10553,7 +10553,7 @@ static void ggml_compute_forward_add( case GGML_TYPE_IQ2_BN: case GGML_TYPE_IQ2_BN_R4: case GGML_TYPE_IQ4_NL: - case GGML_TYPE_IQ4_NL_X4: + case GGML_TYPE_IQ4_NL_R4: case GGML_TYPE_IQ4_XS_R4: case GGML_TYPE_Q4_0_R4: case GGML_TYPE_Q5_0_R4: @@ -11002,7 +11002,7 @@ static void ggml_compute_forward_add1( case GGML_TYPE_IQ2_BN: case GGML_TYPE_IQ2_BN_R4: case GGML_TYPE_IQ4_NL: - case GGML_TYPE_IQ4_NL_X4: + case GGML_TYPE_IQ4_NL_R4: case GGML_TYPE_IQ4_XS_R4: case GGML_TYPE_Q4_0_R4: case GGML_TYPE_Q5_0_R4: @@ -11148,7 +11148,7 @@ static void ggml_compute_forward_acc( case GGML_TYPE_IQ2_BN: case GGML_TYPE_IQ2_BN_R4: case GGML_TYPE_IQ4_NL: - case GGML_TYPE_IQ4_NL_X4: + case GGML_TYPE_IQ4_NL_R4: case GGML_TYPE_IQ4_XS_R4: case GGML_TYPE_Q4_0_R4: case GGML_TYPE_Q5_0_R4: @@ -14340,7 +14340,7 @@ static void ggml_compute_forward_out_prod( case GGML_TYPE_IQ2_BN: case GGML_TYPE_IQ2_BN_R4: case GGML_TYPE_IQ4_NL: - case GGML_TYPE_IQ4_NL_X4: + case GGML_TYPE_IQ4_NL_R4: case GGML_TYPE_IQ4_XS_R4: case GGML_TYPE_Q4_0_R4: case GGML_TYPE_Q5_0_R4: @@ -14726,7 +14726,7 @@ static void ggml_compute_forward_set( case GGML_TYPE_IQ2_BN: case GGML_TYPE_IQ2_BN_R4: case GGML_TYPE_IQ4_NL: - case GGML_TYPE_IQ4_NL_X4: + case GGML_TYPE_IQ4_NL_R4: case GGML_TYPE_IQ4_XS_R4: case GGML_TYPE_Q4_0_R4: case GGML_TYPE_Q5_0_R4: @@ -15006,7 +15006,7 @@ static void ggml_compute_forward_get_rows( case GGML_TYPE_IQ2_BN: case GGML_TYPE_IQ2_BN_R4: case GGML_TYPE_IQ4_NL: - case GGML_TYPE_IQ4_NL_X4: + case GGML_TYPE_IQ4_NL_R4: case GGML_TYPE_IQ4_XS_R4: case GGML_TYPE_Q4_0_R4: case GGML_TYPE_Q5_0_R4: @@ -15613,7 +15613,7 @@ static void ggml_compute_forward_clamp( case GGML_TYPE_IQ2_BN: case GGML_TYPE_IQ2_BN_R4: case GGML_TYPE_IQ4_NL: - case GGML_TYPE_IQ4_NL_X4: + case GGML_TYPE_IQ4_NL_R4: case GGML_TYPE_IQ4_XS_R4: case GGML_TYPE_Q4_0_R4: case GGML_TYPE_Q5_0_R4: @@ -22449,7 +22449,7 @@ size_t ggml_quantize_chunk( case GGML_TYPE_IQ2_BN: result = quantize_iq2_bn (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; 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_X4: result = quantize_iq4_nl_x4(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_Q5_0_R4: result = quantize_q5_0_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 ba6ad15d..69c0b9a4 100644 --- a/ggml/src/iqk/iqk_mul_mat.cpp +++ b/ggml/src/iqk/iqk_mul_mat.cpp @@ -167,7 +167,7 @@ struct MulMat { case GGML_TYPE_Q5_0_R4: case GGML_TYPE_Q6_0_R4: case GGML_TYPE_Q8_0_R4: - case GGML_TYPE_IQ4_NL_X4: + case GGML_TYPE_IQ4_NL_R4: case GGML_TYPE_IQ2_BN_R4: return 4; default: return 1; } @@ -2340,7 +2340,7 @@ static void mul_mat_iq2_bn_r4_q8_k16(int n, const void * vx, size_t bx, const Da #ifdef HAVE_FANCY_SIMD template <int nrc_y> -static void mul_mat_iq4_nl_x4_q8_1(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { +static void mul_mat_iq4_nl_r4_q8_1(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { GGML_ASSERT(nrc_x%8 == 0); Q8<nrc_y, block_q8_1_x4> q8(info); auto m4 = _mm512_set1_epi8(0xf); @@ -2350,8 +2350,8 @@ static void mul_mat_iq4_nl_x4_q8_1(int n, const void * vx, size_t bx, const Data __m512 acc[2*nrc_y] = {}; __m512i qx[4]; for (int ix = 0; ix < nrc_x; ix += 8) { - const block_iq4_nl_x4 * iq4l = (const block_iq4_nl_x4 *)((const char *)vx + (ix+0)*bx); - const block_iq4_nl_x4 * iq4h = (const block_iq4_nl_x4 *)((const char *)vx + (ix+4)*bx); + const block_iq4_nl_r4 * iq4l = (const block_iq4_nl_r4 *)((const char *)vx + (ix+0)*bx); + const block_iq4_nl_r4 * iq4h = (const block_iq4_nl_r4 *)((const char *)vx + (ix+4)*bx); for (int ib4 = 0; ib4 < nb/4; ++ib4) { for (int k = 0; k < 4; ++k) { auto scales128 = _mm_cvtph_ps(_mm_loadl_epi64((const __m128i *)iq4l[4*ib4+k].d)); @@ -2394,7 +2394,7 @@ static void mul_mat_iq4_nl_x4_q8_1(int n, const void * vx, size_t bx, const Data } #else template <int nrc_y> -static void mul_mat_iq4_nl_x4_q8_1(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { +static void mul_mat_iq4_nl_r4_q8_1(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { GGML_ASSERT(nrc_x%4 == 0); Q8<nrc_y, block_q8_1_x4> q8(info); auto m4 = _mm256_set1_epi8(0xf); @@ -2405,7 +2405,7 @@ static void mul_mat_iq4_nl_x4_q8_1(int n, const void * vx, size_t bx, const Data __m256 acc[nrc_y] = {}; //__m256 acc[2*nrc_y] = {}; for (int ix = 0; ix < nrc_x; ix += 4) { - const block_iq4_nl_x4 * iq4 = (const block_iq4_nl_x4 *)((const char *)vx + ix*bx); + const block_iq4_nl_r4 * iq4 = (const block_iq4_nl_r4 *)((const char *)vx + ix*bx); for (int ib4 = 0; ib4 < nb/4; ++ib4) { for (int k = 0; k < 4; ++k) { auto scales128 = _mm_cvtph_ps(_mm_loadl_epi64((const __m128i *)iq4[4*ib4+k].d)); @@ -2455,8 +2455,8 @@ static void mul_mat_q4_0_r4_q8_1(int n, const void * vx, size_t bx, const DataIn __m512 acc[2*nrc_y] = {}; __m512i qx[4]; for (int ix = 0; ix < nrc_x; ix += 8) { - const block_iq4_nl_x4 * iq4l = (const block_iq4_nl_x4 *)((const char *)vx + (ix+0)*bx); - const block_iq4_nl_x4 * iq4h = (const block_iq4_nl_x4 *)((const char *)vx + (ix+4)*bx); + const block_iq4_nl_r4 * iq4l = (const block_iq4_nl_r4 *)((const char *)vx + (ix+0)*bx); + const block_iq4_nl_r4 * iq4h = (const block_iq4_nl_r4 *)((const char *)vx + (ix+4)*bx); for (int ib4 = 0; ib4 < nb/4; ++ib4) { for (int k = 0; k < 4; ++k) { auto scales128 = _mm_cvtph_ps(_mm_loadl_epi64((const __m128i *)iq4l[4*ib4+k].d)); @@ -2508,7 +2508,7 @@ static void mul_mat_q4_0_r4_q8_1(int n, const void * vx, size_t bx, const DataIn GGML_ASSERT(nb%4 == 0); __m256 acc[nrc_y] = {}; for (int ix = 0; ix < nrc_x; ix += 4) { - const block_iq4_nl_x4 * iq4 = (const block_iq4_nl_x4 *)((const char *)vx + ix*bx); + const block_iq4_nl_r4 * iq4 = (const block_iq4_nl_r4 *)((const char *)vx + ix*bx); for (int ib4 = 0; ib4 < nb/4; ++ib4) { for (int k = 0; k < 4; ++k) { auto scales128 = _mm_cvtph_ps(_mm_loadl_epi64((const __m128i *)iq4[4*ib4+k].d)); @@ -5052,16 +5052,16 @@ bool MulMat::prepare(int typeA, int typeB, int ne00, MulMat& mm, int Ny) { MulMat::set_functions<IQ4_NL_Unpacker>(mm); expected_typeB = GGML_TYPE_Q8_1; break; - case GGML_TYPE_IQ4_NL_X4: + case GGML_TYPE_IQ4_NL_R4: assert (ne00 % QK4_NL == 0); - mm.funcs[0] = mul_mat_iq4_nl_x4_q8_1<1>; - mm.funcs[1] = mul_mat_iq4_nl_x4_q8_1<2>; - mm.funcs[2] = mul_mat_iq4_nl_x4_q8_1<3>; - mm.funcs[3] = mul_mat_iq4_nl_x4_q8_1<4>; - mm.funcs[4] = mul_mat_iq4_nl_x4_q8_1<5>; - mm.funcs[5] = mul_mat_iq4_nl_x4_q8_1<6>; - mm.funcs[6] = mul_mat_iq4_nl_x4_q8_1<7>; - mm.funcs[7] = mul_mat_iq4_nl_x4_q8_1<8>; + mm.funcs[0] = mul_mat_iq4_nl_r4_q8_1<1>; + mm.funcs[1] = mul_mat_iq4_nl_r4_q8_1<2>; + mm.funcs[2] = mul_mat_iq4_nl_r4_q8_1<3>; + mm.funcs[3] = mul_mat_iq4_nl_r4_q8_1<4>; + mm.funcs[4] = mul_mat_iq4_nl_r4_q8_1<5>; + mm.funcs[5] = mul_mat_iq4_nl_r4_q8_1<6>; + mm.funcs[6] = mul_mat_iq4_nl_r4_q8_1<7>; + mm.funcs[7] = mul_mat_iq4_nl_r4_q8_1<8>; expected_typeB = GGML_TYPE_Q8_1; break; case GGML_TYPE_IQ4_XS_R4: @@ -7734,7 +7734,7 @@ void mul_mat_iq4_xs_r4_q8_k(int n, const void * vx, size_t bx, const DataInfo& i } } -void mul_mat_iq4_nl_x4_q8_0_1(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { +void mul_mat_iq4_nl_r4_q8_0_1(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) { GGML_ASSERT(nrc_x%4 == 0); Q8<1, block_q8_0_x4> q8(info); auto m4 = vdupq_n_u8(0xf); @@ -7744,7 +7744,7 @@ void mul_mat_iq4_nl_x4_q8_0_1(int n, const void * vx, size_t bx, const DataInfo& int8x16_t qx[8]; for (int ix = 0; ix < nrc_x; ix += 4) { auto acc = vdupq_n_f32(0.f); - const block_iq4_nl_x4 * iq4 = (const block_iq4_nl_x4 *)((const char *)vx + ix*bx); + const block_iq4_nl_r4 * iq4 = (const block_iq4_nl_r4 *)((const char *)vx + ix*bx); for (int ib4 = 0; ib4 < nb/4; ++ib4) { auto y1 = vld1q_s8_x4(q8.y[0][ib4].qs); auto y2 = vld1q_s8_x4(q8.y[0][ib4].qs+64); @@ -7812,7 +7812,7 @@ void mul_mat_qx_r4_q8_0(int n, const void * vx, size_t bx, const DataInfo& info, struct IQ4_NL_R4_Dequantizer { IQ4_NL_R4_Dequantizer(const void * vx, size_t bx) : cx((const char *)vx), bx(bx), values(vld1q_s8(iq4k_values)) {} - inline void new_row(int ix) { iq4 = (const block_iq4_nl_x4 *)(cx + ix*bx); } + inline void new_row(int ix) { iq4 = (const block_iq4_nl_r4 *)(cx + ix*bx); } inline float32x4_t prepare(int ib4, int k, int8x16_t * qx) const { auto scales = vcvt_f32_f16(vld1_f16((const float16_t *)iq4[4*ib4+k].d)); auto bits = vld1q_u8_x4(iq4[4*ib4+k].qs); @@ -7825,14 +7825,14 @@ struct IQ4_NL_R4_Dequantizer { const char * cx; const size_t bx; - const block_iq4_nl_x4 * iq4; + const block_iq4_nl_r4 * iq4; const uint8x16_t m4 = vdupq_n_u8(0x0f); const int8x16_t values; }; struct Q4_0_R4_Dequantizer { Q4_0_R4_Dequantizer(const void * vx, size_t bx) : cx((const char *)vx), bx(bx) {} - inline void new_row(int ix) { iq4 = (const block_iq4_nl_x4 *)(cx + ix*bx); } + inline void new_row(int ix) { iq4 = (const block_iq4_nl_r4 *)(cx + ix*bx); } inline float32x4_t prepare(int ib4, int k, int8x16_t * qx) const { auto scales = vcvt_f32_f16(vld1_f16((const float16_t *)iq4[4*ib4+k].d)); auto bits = vld1q_u8_x4(iq4[4*ib4+k].qs); @@ -7853,7 +7853,7 @@ struct Q4_0_R4_Dequantizer { const char * cx; const size_t bx; - const block_iq4_nl_x4 * iq4; + const block_iq4_nl_r4 * iq4; const uint8x16_t m4 = vdupq_n_u8(0xf0); const uint8x16_t m88 = vdupq_n_u8(0x88); const float32x4_t norm = vdupq_n_f32(1.f/16); @@ -8123,7 +8123,7 @@ bool MulMat::prepare(int typeA, int typeB, int ne00, MulMat& m, int /*Ny*/) { MulMat::set_functions<DequantizerIQ4NL>(m); expected_Btype = GGML_TYPE_Q8_0; break; - case GGML_TYPE_IQ4_NL_X4: + case GGML_TYPE_IQ4_NL_R4: SET_MUL_MAT_FUNCTIONS_T(m, mul_mat_qx_r4_q8_0, IQ4_NL_R4_Dequantizer); expected_Btype = GGML_TYPE_Q8_0; break; diff --git a/ggml/src/iqk/iqk_quantize.cpp b/ggml/src/iqk/iqk_quantize.cpp index 32fe92ef..0ea608bc 100644 --- a/ggml/src/iqk/iqk_quantize.cpp +++ b/ggml/src/iqk/iqk_quantize.cpp @@ -3233,19 +3233,19 @@ void vec_dot_iq4_kss_q8_k(int n, float * s, size_t bs, const void * vx, size_t b } // -// ========================================= iq4_nl_x4 +// ========================================= iq4_nl_r4 // -void quantize_row_iq4_nl_x4_ref(const float * x, block_iq4_nl_x4 * y, int64_t k) { +void quantize_row_iq4_nl_r4_ref(const float * x, block_iq4_nl_r4 * y, int64_t k) { // we assume we are called with 4 rows - quantize_iq4_nl_x4(x, (void *)y, 4, k/4, nullptr); + quantize_iq4_nl_r4(x, (void *)y, 4, k/4, nullptr); } -void quantize_row_iq4_nl_x4(const float * x, void * y, int64_t k) { +void quantize_row_iq4_nl_r4(const float * x, void * y, int64_t k) { // we assume we are called with 4 rows - quantize_iq4_nl_x4(x, y, 4, k/4, nullptr); + quantize_iq4_nl_r4(x, y, 4, k/4, nullptr); } -static void repack_iq4_nl(int nrows, int n_per_row, const block_iq4_nl * x, block_iq4_nl_x4 * y) { +static void repack_iq4_nl(int nrows, int n_per_row, const block_iq4_nl * x, block_iq4_nl_r4 * y) { GGML_ASSERT(nrows%4 == 0); GGML_ASSERT(n_per_row%QK4_NL == 0); int nblock = n_per_row/QK4_NL; @@ -3266,21 +3266,21 @@ static void repack_iq4_nl(int nrows, int n_per_row, const block_iq4_nl * x, bloc } } -size_t quantize_iq4_nl_x4(const float * src, void * dst, int64_t nrows, int64_t n_per_row, const float * imatrix) { +size_t quantize_iq4_nl_r4(const float * src, void * dst, int64_t nrows, int64_t n_per_row, const float * imatrix) { GGML_ASSERT(nrows%4 == 0); auto row_size_nl = ggml_row_size(GGML_TYPE_IQ4_NL, n_per_row); std::vector<char> qtmp(4*row_size_nl); char * qrow = (char *)dst; for (int row = 0; row < nrows; row += 4) { quantize_iq4_nl(src, qtmp.data(), 4, n_per_row, imatrix); - repack_iq4_nl(4, n_per_row, (const block_iq4_nl *)qtmp.data(), (block_iq4_nl_x4 *)qrow); + repack_iq4_nl(4, n_per_row, (const block_iq4_nl *)qtmp.data(), (block_iq4_nl_r4 *)qrow); src += 4*n_per_row; qrow += 4*row_size_nl; } return nrows*row_size_nl; } -void dequantize_row_iq4_nl_x4(const block_iq4_nl_x4 * x, float * y, int64_t k) { +void dequantize_row_iq4_nl_r4(const block_iq4_nl_r4 * x, float * y, int64_t k) { // we assume we are called with 4 rows int n_per_row = k/4; int nb = n_per_row/QK4_NL; @@ -3303,9 +3303,9 @@ void dequantize_row_iq4_nl_x4(const block_iq4_nl_x4 * x, float * y, int64_t k) { } } -void vec_dot_iq4_nl_x4_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_iq4_nl_r4_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_IQ4_NL_X4, vx, 0, GGML_TYPE_Q8_0, vy, 0, s, 0, 0, 1)) { + if (iqk_mul_mat(1, 1, n, GGML_TYPE_IQ4_NL_R4, vx, 0, GGML_TYPE_Q8_0, vy, 0, s, 0, 0, 1)) { return; } #endif @@ -3319,7 +3319,7 @@ void vec_dot_iq4_nl_x4_q8_0(int n, float * s, size_t bs, const void * vx, size_t // // ========================================= q4_0_r4 // -void quantize_row_q4_0_r4_ref(const float * x, block_iq4_nl_x4 * y, int64_t k) { +void quantize_row_q4_0_r4_ref(const float * x, block_iq4_nl_r4 * y, int64_t k) { // we assume we are called with 4 rows quantize_q4_0_r4(x, (void *)y, 4, k/4, nullptr); } @@ -3329,7 +3329,7 @@ void quantize_row_q4_0_r4(const float * x, void * y, int64_t k) { quantize_q4_0_r4(x, y, 4, k/4, nullptr); } -static void repack_q4_0(int nrows, int n_per_row, const block_q4_0 * x, block_iq4_nl_x4 * y) { +static void repack_q4_0(int nrows, int n_per_row, const block_q4_0 * x, block_iq4_nl_r4 * y) { GGML_ASSERT(nrows%4 == 0); GGML_ASSERT(n_per_row%QK4_NL == 0); int nblock = n_per_row/QK4_NL; @@ -3369,14 +3369,14 @@ size_t quantize_q4_0_r4(const float * src, void * dst, int64_t nrows, int64_t n_ char * qrow = (char *)dst; for (int row = 0; row < nrows; row += 4) { quantize_q4_0(src, qtmp.data(), 4, n_per_row, imatrix); - repack_iq4_nl(4, n_per_row, (const block_iq4_nl *)qtmp.data(), (block_iq4_nl_x4 *)qrow); + repack_iq4_nl(4, n_per_row, (const block_iq4_nl *)qtmp.data(), (block_iq4_nl_r4 *)qrow); src += 4*n_per_row; qrow += 4*row_size_nl; } return nrows*row_size_nl; } -void dequantize_row_q4_0_r4(const block_iq4_nl_x4 * x, float * y, int64_t k) { +void dequantize_row_q4_0_r4(const block_iq4_nl_r4 * x, float * y, int64_t k) { // we assume we are called with 4 rows int n_per_row = k/4; int nb = n_per_row/QK4_0; diff --git a/ggml/src/iqk/iqk_quantize.h b/ggml/src/iqk/iqk_quantize.h index fb436e57..f5bbd3dc 100644 --- a/ggml/src/iqk/iqk_quantize.h +++ b/ggml/src/iqk/iqk_quantize.h @@ -63,16 +63,16 @@ void vec_dot_iq2_ks_q8_k(int n, float * GGML_RESTRICT s, size_t bs, const void void iqk_quantize_row_q8_K(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k); -void quantize_row_iq4_nl_x4_ref(const float * GGML_RESTRICT x, block_iq4_nl_x4 * GGML_RESTRICT y, int64_t k); -void quantize_row_iq4_nl_x4(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); -size_t quantize_iq4_nl_x4(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); -void dequantize_row_iq4_nl_x4(const block_iq4_nl_x4 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); -void vec_dot_iq4_nl_x4_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_nl_r4_ref(const float * GGML_RESTRICT x, block_iq4_nl_r4 * GGML_RESTRICT y, int64_t k); +void quantize_row_iq4_nl_r4(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +size_t quantize_iq4_nl_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_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_x4 * GGML_RESTRICT y, int64_t k); +void quantize_row_q4_0_r4_ref(const float * GGML_RESTRICT x, block_iq4_nl_r4 * 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_x4 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +void dequantize_row_q4_0_r4(const block_iq4_nl_r4 * 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_x4 * GGML_RESTRICT y, int64_t k); diff --git a/include/llama.h b/include/llama.h index 243fddaa..9eec6a43 100644 --- a/include/llama.h +++ b/include/llama.h @@ -183,7 +183,7 @@ extern "C" { LLAMA_FTYPE_MOSTLY_Q4_0_R4 = 202, // except 1d tensors LLAMA_FTYPE_MOSTLY_Q8_0_R4 = 207, // except 1d tensors LLAMA_FTYPE_MOSTLY_Q5_0_R4 = 208, // except 1d tensors - LLAMA_FTYPE_MOSTLY_IQ4_NL_X4 = 225, // except 1d tensors + LLAMA_FTYPE_MOSTLY_IQ4_NL_R4 = 225, // except 1d tensors LLAMA_FTYPE_MOSTLY_IQ4_XS_R4 = 230, // except 1d tensors LLAMA_FTYPE_MOSTLY_Q6_0_R4 = 235, // except 1d tensors LLAMA_FTYPE_MOSTLY_IQ2_BN_R4 = 237, // except 1d tensors diff --git a/src/llama.cpp b/src/llama.cpp index 0e1aadbd..4b25d4ec 100644 --- a/src/llama.cpp +++ b/src/llama.cpp @@ -3850,7 +3850,7 @@ struct llama_model_loader { case GGML_TYPE_IQ2_BN: ftype = LLAMA_FTYPE_MOSTLY_IQ2_BN; break; case GGML_TYPE_IQ2_BN_R4:ftype = LLAMA_FTYPE_MOSTLY_IQ2_BN_R4;break; case GGML_TYPE_IQ4_NL: ftype = LLAMA_FTYPE_MOSTLY_IQ4_NL; break; - case GGML_TYPE_IQ4_NL_X4:ftype = LLAMA_FTYPE_MOSTLY_IQ4_NL_X4;break; + case GGML_TYPE_IQ4_NL_R4:ftype = LLAMA_FTYPE_MOSTLY_IQ4_NL_R4;break; case GGML_TYPE_IQ4_XS_R4:ftype = LLAMA_FTYPE_MOSTLY_IQ4_XS_R4;break; case GGML_TYPE_Q4_0_R4: ftype = LLAMA_FTYPE_MOSTLY_Q4_0_R4; break; case GGML_TYPE_Q5_0_R4: ftype = LLAMA_FTYPE_MOSTLY_Q5_0_R4; break; @@ -4560,7 +4560,7 @@ static std::string llama_model_ftype_name(llama_ftype ftype) { case LLAMA_FTYPE_MOSTLY_IQ1_S: return "IQ1_S - 1.5625 bpw"; case LLAMA_FTYPE_MOSTLY_IQ1_M: return "IQ1_M - 1.75 bpw"; case LLAMA_FTYPE_MOSTLY_IQ4_NL: return "IQ4_NL - 4.5 bpw"; - case LLAMA_FTYPE_MOSTLY_IQ4_NL_X4:return "IQ4_NL_X4 - 4.5 bpw"; + case LLAMA_FTYPE_MOSTLY_IQ4_NL_R4:return "IQ4_NL_R4 - 4.5 bpw"; case LLAMA_FTYPE_MOSTLY_IQ4_XS_R4:return "IQ4_XS_R4 - 4.25 bpw"; case LLAMA_FTYPE_MOSTLY_Q4_0_R4: return "Q4_0_R4 - 4.5 bpw"; case LLAMA_FTYPE_MOSTLY_Q5_0_R4: return "Q5_0_R4 - 5.5 bpw"; @@ -15780,7 +15780,7 @@ static ggml_type llama_tensor_get_type(quantize_state_internal & qs, ggml_type n new_type == GGML_TYPE_Q4_0_8_8) { new_type = GGML_TYPE_Q4_0; } - else if (new_type == GGML_TYPE_IQ4_NL_X4) { + else if (new_type == GGML_TYPE_IQ4_NL_R4) { new_type = GGML_TYPE_IQ4_NL; } else if (new_type == GGML_TYPE_IQ4_XS_R4) { @@ -15860,7 +15860,7 @@ static ggml_type llama_tensor_get_type(quantize_state_internal & qs, ggml_type n } else if (ftype == LLAMA_FTYPE_MOSTLY_Q3_K_L) new_type = GGML_TYPE_Q5_K; else if ((ftype == LLAMA_FTYPE_MOSTLY_IQ4_NL || ftype == LLAMA_FTYPE_MOSTLY_IQ4_XS || - ftype == LLAMA_FTYPE_MOSTLY_IQ4_NL_X4 || ftype == LLAMA_FTYPE_MOSTLY_IQ4_XS_R4 || + ftype == LLAMA_FTYPE_MOSTLY_IQ4_NL_R4 || ftype == LLAMA_FTYPE_MOSTLY_IQ4_XS_R4 || ftype == LLAMA_FTYPE_MOSTLY_IQ4_KS || ftype == LLAMA_FTYPE_MOSTLY_IQ4_KSS) && qs.model.hparams.n_gqa() >= 2) { new_type = GGML_TYPE_IQ5_K; } @@ -15890,7 +15890,7 @@ static ggml_type llama_tensor_get_type(quantize_state_internal & qs, ggml_type n else if (new_type == GGML_TYPE_Q3_K || new_type == GGML_TYPE_IQ3_S ) new_type = GGML_TYPE_Q4_K; else if (new_type == GGML_TYPE_Q4_K || new_type == GGML_TYPE_IQ4_XS) new_type = GGML_TYPE_Q5_K; else if (new_type == GGML_TYPE_IQ4_NL) new_type = GGML_TYPE_Q5_K; - else if (new_type == GGML_TYPE_IQ4_NL_X4) new_type = GGML_TYPE_Q5_K; + else if (new_type == GGML_TYPE_IQ4_NL_R4) new_type = GGML_TYPE_Q5_K; else if (new_type == GGML_TYPE_IQ4_XS_R4) new_type = GGML_TYPE_Q5_K; else if (new_type == GGML_TYPE_Q5_K) new_type = GGML_TYPE_Q6_K; } @@ -15957,7 +15957,7 @@ static ggml_type llama_tensor_get_type(quantize_state_internal & qs, ggml_type n else if (i_layer < n_layer/8 && !qs.has_imatrix && (ftype == LLAMA_FTYPE_MOSTLY_IQ4_NL || ftype == LLAMA_FTYPE_MOSTLY_IQ4_XS || ftype == LLAMA_FTYPE_MOSTLY_IQ4_KS || ftype == LLAMA_FTYPE_MOSTLY_IQ4_KSS || - ftype == LLAMA_FTYPE_MOSTLY_IQ4_NL_X4 || ftype == LLAMA_FTYPE_MOSTLY_IQ4_XS_R4)) { + ftype == LLAMA_FTYPE_MOSTLY_IQ4_NL_R4 || ftype == LLAMA_FTYPE_MOSTLY_IQ4_XS_R4)) { new_type = GGML_TYPE_Q5_K; } else if (ftype == LLAMA_FTYPE_MOSTLY_Q5_K_M && use_more_bits(i_layer, n_layer)) new_type = GGML_TYPE_Q6_K; @@ -15972,7 +15972,7 @@ static ggml_type llama_tensor_get_type(quantize_state_internal & qs, ggml_type n new_type = ftype == LLAMA_FTYPE_MOSTLY_Q4_0 ? GGML_TYPE_Q4_1 : GGML_TYPE_Q5_1; } else if (ftype == LLAMA_FTYPE_MOSTLY_Q4_0_R4 && qs.has_imatrix && i_layer < n_layer/8) { - new_type = GGML_TYPE_IQ4_NL_X4; + new_type = GGML_TYPE_IQ4_NL_R4; } ++qs.i_ffn_down; } else if (name.find("attn_output.weight") != std::string::npos) { @@ -15984,7 +15984,7 @@ static ggml_type llama_tensor_get_type(quantize_state_internal & qs, ggml_type n ftype == LLAMA_FTYPE_MOSTLY_Q4_K_S || ftype == LLAMA_FTYPE_MOSTLY_Q4_K_M || ftype == LLAMA_FTYPE_MOSTLY_IQ3_S || ftype == LLAMA_FTYPE_MOSTLY_IQ3_M || ftype == LLAMA_FTYPE_MOSTLY_IQ4_XS || ftype == LLAMA_FTYPE_MOSTLY_IQ4_K || ftype == LLAMA_FTYPE_MOSTLY_IQ2_K || ftype == LLAMA_FTYPE_MOSTLY_IQ3_K || - ftype == LLAMA_FTYPE_MOSTLY_IQ4_NL_X4 || ftype == LLAMA_FTYPE_MOSTLY_IQ4_XS_R4) { + ftype == LLAMA_FTYPE_MOSTLY_IQ4_NL_R4 || ftype == LLAMA_FTYPE_MOSTLY_IQ4_XS_R4) { new_type = GGML_TYPE_Q5_K; } } else { @@ -16194,7 +16194,7 @@ static void llama_model_quantize_internal(const std::string & fname_inp, const s case LLAMA_FTYPE_MOSTLY_IQ2_BN: default_type = GGML_TYPE_IQ2_BN; break; case LLAMA_FTYPE_MOSTLY_IQ2_BN_R4:default_type = GGML_TYPE_IQ2_BN_R4;break; case LLAMA_FTYPE_MOSTLY_IQ4_NL: default_type = GGML_TYPE_IQ4_NL; break; - case LLAMA_FTYPE_MOSTLY_IQ4_NL_X4:default_type = GGML_TYPE_IQ4_NL_X4;break; + case LLAMA_FTYPE_MOSTLY_IQ4_NL_R4:default_type = GGML_TYPE_IQ4_NL_R4;break; case LLAMA_FTYPE_MOSTLY_IQ4_XS_R4:default_type = GGML_TYPE_IQ4_XS_R4;break; case LLAMA_FTYPE_MOSTLY_Q4_0_R4: default_type = GGML_TYPE_Q4_0_R4; break; case LLAMA_FTYPE_MOSTLY_Q5_0_R4: default_type = GGML_TYPE_Q5_0_R4; break; @@ -16557,7 +16557,7 @@ static void llama_model_quantize_internal(const std::string & fname_inp, const s if (new_type == GGML_TYPE_Q4_0_8_8) chunk_size_multiplier = 8; else if (new_type == GGML_TYPE_Q4_0_4_4 || new_type == GGML_TYPE_Q4_0_4_8) chunk_size_multiplier = 4; } - else if (new_type == GGML_TYPE_IQ4_NL_X4) { + else if (new_type == GGML_TYPE_IQ4_NL_R4) { if (tensor->ne[1] % 4 != 0) new_type = GGML_TYPE_IQ4_NL; else chunk_size_multiplier = 4; } |