nicolas
/
ik_llama_opt
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

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 <iwan.kawrakow@gmail.com>
This commit is contained in:
Kawrakow 2025-02-06 18:45:28 +02:00 committed by GitHub
parent 7c94c3da56
commit 8049ffcbc8
10 changed files with 196 additions and 196 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
examples/quantize/quantize.cpp
View File

@ -51,11 +51,11 @@ static const std::vector<struct quant_option> QUANT_OPTIONS = {
{ "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_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", },
{ "IQ4_XS_R8",LLAMA_FTYPE_MOSTLY_IQ4_XS_R8," 4.25 bpw non-linear quantization", },
{ "Q4_0_R8", LLAMA_FTYPE_MOSTLY_Q4_0_R8, " 4.50 bpw quantization", },
{ "Q5_0_R4", LLAMA_FTYPE_MOSTLY_Q5_0_R4, " 5.50 bpw quantization", },
{ "Q6_0_R4", LLAMA_FTYPE_MOSTLY_Q6_0_R4, " 6.50 bpw quantization", },
{ "Q8_0_R4", LLAMA_FTYPE_MOSTLY_Q8_0_R4, " 8.50 bpw quantization", },
{ "Q8_0_R8", LLAMA_FTYPE_MOSTLY_Q8_0_R8, " 8.50 bpw quantization", },
{ "IQ4_XS", LLAMA_FTYPE_MOSTLY_IQ4_XS, " 4.25 bpw non-linear quantization", },
{ "IQ4_KS", LLAMA_FTYPE_MOSTLY_IQ4_KS, " 4.25 bpw non-linear quantization", },
{ "IQ4_KS_R4",LLAMA_FTYPE_MOSTLY_IQ4_KS_R4,"IQ4_KS repacked", },

12
ggml/include/ggml.h
View File

@ -416,9 +416,9 @@ extern "C" {
GGML_TYPE_Q8_K32 = 148,
GGML_TYPE_Q8_KR8 = 149,
GGML_TYPE_Q4_0_R4 = 202,
GGML_TYPE_Q4_0_R8 = 202,
GGML_TYPE_Q5_0_R4 = 206,
GGML_TYPE_Q8_0_R4 = 208,
GGML_TYPE_Q8_0_R8 = 208,
GGML_TYPE_Q2_K_R4 = 210,
GGML_TYPE_Q3_K_R4 = 211,
GGML_TYPE_Q4_K_R4 = 212,
@ -431,7 +431,7 @@ extern "C" {
GGML_TYPE_IQ4_NL_R4 = 220,
GGML_TYPE_IQ3_S_R4 = 221,
GGML_TYPE_IQ2_S_R4 = 222,
GGML_TYPE_IQ4_XS_R4 = 223,
GGML_TYPE_IQ4_XS_R8 = 223,
GGML_TYPE_IQ1_M_R4 = 229,
GGML_TYPE_BF16_R16 = 230,
GGML_TYPE_Q6_0_R4 = 233,
@ -501,8 +501,8 @@ extern "C" {
GGML_FTYPE_MOSTLY_IQ2_KS = 138, // except 1d tensors
GGML_FTYPE_MOSTLY_IQ4_KSS = 139, // except 1d tensors
//
GGML_FTYPE_MOSTLY_Q4_0_R4 = 202, // except 1d tensors
GGML_FTYPE_MOSTLY_Q8_0_R4 = 207, // except 1d tensors
GGML_FTYPE_MOSTLY_Q4_0_R8 = 202, // except 1d tensors
GGML_FTYPE_MOSTLY_Q8_0_R8 = 207, // except 1d tensors
GGML_FTYPE_MOSTLY_Q5_0_R4 = 208, // except 1d tensors
GGML_FTYPE_MOSTLY_Q2_K_R4 = 210, // except 1d tensors
GGML_FTYPE_MOSTLY_Q3_K_R4 = 211, // except 1d tensors
@ -516,7 +516,7 @@ extern "C" {
GGML_FTYPE_MOSTLY_IQ4_NL_R4 = 219, // except 1d tensors
GGML_FTYPE_MOSTLY_IQ3_S_R4 = 220, // except 1d tensors
GGML_FTYPE_MOSTLY_IQ2_S_R4 = 221, // except 1d tensors
GGML_FTYPE_MOSTLY_IQ4_XS_R4 = 222, // except 1d tensors
GGML_FTYPE_MOSTLY_IQ4_XS_R8 = 222, // except 1d tensors
GGML_FTYPE_MOSTLY_IQ1_M_R4 = 223, // except 1d tensors
GGML_FTYPE_MOSTLY_BF16_R16 = 224, // except 1d tensors
GGML_FTYPE_MOSTLY_Q6_0_R4 = 227, // except 1d tensors

4
ggml/src/ggml-common.h
View File

@ -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];

6
ggml/src/ggml-quants.c
View File

@ -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;

90
ggml/src/ggml.c
View File

@ -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;

120
ggml/src/iqk/iqk_mul_mat.cpp
View File

@ -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 <int nrc_y>
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<nrc_y, block_q8_1_x4> 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 <int nrc_y>
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 <int nrc_y>
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<nrc_y>(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<nrc_y>(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 <int nrc_y>
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<nrc_y, block_q8_1_x4> 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 <int nrc_y>
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<nrc_y, block_q8_1_x4> 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 <int nrc_y>
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<nrc_y, block_q8_K> 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 <int nrc_y>
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<nrc_y>(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<nrc_y>(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<nrc_y, block_q8_K> 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 <int nrc_y>
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<nrc_y>(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<nrc_y>(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 <int nrc_y>
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<nrc_y, block_q8_K> 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 <int nrc_y>
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<nrc_y, block_q8_0_x4> 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<KHelper, HelperQ80R4<D, k_step>>) {
#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<KHelper, HelperQ41<D, k_step>>) {

72
ggml/src/iqk/iqk_quantize.cpp
View File

@ -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<char> 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<char> 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<char> 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<ggml_type, Modify> 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<ggml_bf16_t>}},

30
ggml/src/iqk/iqk_quantize.h
View File

@ -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_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_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_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);

6
include/llama.h
View File

@ -181,8 +181,8 @@ extern "C" {
LLAMA_FTYPE_MOSTLY_IQ2_KS = 147, // except 1d tensors
LLAMA_FTYPE_MOSTLY_IQ4_KSS = 148, // except 1d tensors
//
LLAMA_FTYPE_MOSTLY_Q4_0_R4 = 202, // except 1d tensors
LLAMA_FTYPE_MOSTLY_Q8_0_R4 = 207, // except 1d tensors
LLAMA_FTYPE_MOSTLY_Q4_0_R8 = 202, // except 1d tensors
LLAMA_FTYPE_MOSTLY_Q8_0_R8 = 207, // except 1d tensors
LLAMA_FTYPE_MOSTLY_Q5_0_R4 = 208, // except 1d tensors
LLAMA_FTYPE_MOSTLY_Q2_K_R4 = 210, // except 1d tensors
LLAMA_FTYPE_MOSTLY_Q3_K_R4 = 211, // except 1d tensors
@ -196,7 +196,7 @@ extern "C" {
LLAMA_FTYPE_MOSTLY_IQ4_NL_R4 = 225, // except 1d tensors
LLAMA_FTYPE_MOSTLY_IQ3_S_R4 = 226, // except 1d tensors
LLAMA_FTYPE_MOSTLY_IQ2_M_R4 = 229, // except 1d tensors
LLAMA_FTYPE_MOSTLY_IQ4_XS_R4 = 230, // except 1d tensors
LLAMA_FTYPE_MOSTLY_IQ4_XS_R8 = 230, // except 1d tensors
LLAMA_FTYPE_MOSTLY_IQ1_M_R4 = 231, // except 1d tensors
LLAMA_FTYPE_MOSTLY_Q6_0_R4 = 335, // except 1d tensors
LLAMA_FTYPE_MOSTLY_BF16_R16 = 232, // except 1d tensors

46
src/llama.cpp
View File

@ -3962,11 +3962,11 @@ struct llama_model_loader {
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_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_IQ4_XS_R8:ftype = LLAMA_FTYPE_MOSTLY_IQ4_XS_R8;break;
case GGML_TYPE_Q4_0_R8: ftype = LLAMA_FTYPE_MOSTLY_Q4_0_R8; break;
case GGML_TYPE_Q5_0_R4: ftype = LLAMA_FTYPE_MOSTLY_Q5_0_R4; break;
case GGML_TYPE_Q6_0_R4: ftype = LLAMA_FTYPE_MOSTLY_Q6_0_R4; break;
case GGML_TYPE_Q8_0_R4: ftype = LLAMA_FTYPE_MOSTLY_Q8_0_R4; break;
case GGML_TYPE_Q8_0_R8: ftype = LLAMA_FTYPE_MOSTLY_Q8_0_R8; break;
case GGML_TYPE_IQ4_XS: ftype = LLAMA_FTYPE_MOSTLY_IQ4_XS; break;
case GGML_TYPE_IQ4_KS: ftype = LLAMA_FTYPE_MOSTLY_IQ4_KS; break;
case GGML_TYPE_IQ4_KS_R4:ftype = LLAMA_FTYPE_MOSTLY_IQ4_KS_R4; break;
@ -4695,11 +4695,11 @@ static std::string llama_model_ftype_name(llama_ftype ftype) {
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_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_IQ4_XS_R8:return "IQ4_XS_R8 - 4.25 bpw";
case LLAMA_FTYPE_MOSTLY_Q4_0_R8: return "Q4_0_R8 - 4.5 bpw";
case LLAMA_FTYPE_MOSTLY_Q5_0_R4: return "Q5_0_R4 - 5.5 bpw";
case LLAMA_FTYPE_MOSTLY_Q6_0_R4: return "Q6_0_R4 - 6.5 bpw";
case LLAMA_FTYPE_MOSTLY_Q8_0_R4: return "Q8_0_R4 - 8.5 bpw";
case LLAMA_FTYPE_MOSTLY_Q8_0_R8: return "Q8_0_R8 - 8.5 bpw";
case LLAMA_FTYPE_MOSTLY_IQ4_XS: return "IQ4_XS - 4.25 bpw";
case LLAMA_FTYPE_MOSTLY_IQ4_KS: return "IQ4_KS - 4.25 bpw";
case LLAMA_FTYPE_MOSTLY_IQ4_KS_R4:return "IQ4_KS_R4 - 4.25 bpw";
@ -15982,7 +15982,7 @@ static ggml_type llama_tensor_get_type(quantize_state_internal & qs, ggml_type n
ftype == LLAMA_FTYPE_MOSTLY_IQ4_KS || ftype == LLAMA_FTYPE_MOSTLY_IQ4_KSS || ftype == LLAMA_FTYPE_MOSTLY_IQ4_KS_R4) && !qs.has_output) {
new_type = GGML_TYPE_IQ5_K;
}
else if (new_type != GGML_TYPE_Q8_0 && new_type != GGML_TYPE_Q8_0_R4 && new_type != GGML_TYPE_IQ6_K && new_type != GGML_TYPE_Q6_K_R4 &&
else if (new_type != GGML_TYPE_Q8_0 && new_type != GGML_TYPE_Q8_0_R8 && new_type != GGML_TYPE_IQ6_K && new_type != GGML_TYPE_Q6_K_R4 &&
new_type != GGML_TYPE_Q8_K_R8) {
new_type = GGML_TYPE_Q6_K;
}
@ -16016,7 +16016,7 @@ static ggml_type llama_tensor_get_type(quantize_state_internal & qs, ggml_type n
else if (new_type == GGML_TYPE_IQ4_NL_R4) {
new_type = GGML_TYPE_IQ4_NL;
}
else if (new_type == GGML_TYPE_IQ4_XS_R4) {
else if (new_type == GGML_TYPE_IQ4_XS_R8) {
new_type = GGML_TYPE_IQ4_XS;
}
else if (new_type == GGML_TYPE_Q2_K_R4) {
@ -16055,7 +16055,7 @@ static ggml_type llama_tensor_get_type(quantize_state_internal & qs, ggml_type n
else if (new_type == GGML_TYPE_IQ4_KS_R4) {
new_type = GGML_TYPE_IQ4_KS;
}
else if (new_type == GGML_TYPE_Q4_0_R4) {
else if (new_type == GGML_TYPE_Q4_0_R8) {
new_type = GGML_TYPE_Q4_0;
}
else if (new_type == GGML_TYPE_Q5_0_R4) {
@ -16064,7 +16064,7 @@ static ggml_type llama_tensor_get_type(quantize_state_internal & qs, ggml_type n
else if (new_type == GGML_TYPE_Q6_0_R4) {
new_type = GGML_TYPE_Q6_0;
}
else if (new_type == GGML_TYPE_Q8_0_R4) {
else if (new_type == GGML_TYPE_Q8_0_R8) {
new_type = GGML_TYPE_Q8_0;
}
else if (new_type == GGML_TYPE_BF16_R16) {
@ -16188,7 +16188,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_R4 || ftype == LLAMA_FTYPE_MOSTLY_IQ4_XS_R4 ||
ftype == LLAMA_FTYPE_MOSTLY_IQ4_NL_R4 || ftype == LLAMA_FTYPE_MOSTLY_IQ4_XS_R8 ||
ftype == LLAMA_FTYPE_MOSTLY_IQ4_KS || ftype == LLAMA_FTYPE_MOSTLY_IQ4_KSS) && qs.model.hparams.n_gqa() >= 2) {
new_type = GGML_TYPE_IQ5_K;
}
@ -16229,7 +16229,7 @@ static ggml_type llama_tensor_get_type(quantize_state_internal & qs, ggml_type n
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_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_IQ4_XS_R8) new_type = GGML_TYPE_Q5_K;
else if (new_type == GGML_TYPE_Q5_K) new_type = GGML_TYPE_Q6_K;
}
++qs.i_attention_wv;
@ -16306,7 +16306,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_R4 || ftype == LLAMA_FTYPE_MOSTLY_IQ4_XS_R4)) {
ftype == LLAMA_FTYPE_MOSTLY_IQ4_NL_R4 || ftype == LLAMA_FTYPE_MOSTLY_IQ4_XS_R8)) {
new_type = GGML_TYPE_Q5_K;
}
else if (ftype == LLAMA_FTYPE_MOSTLY_IQ4_KS_R4 && i_layer < n_layer/8 && !qs.has_imatrix) {
@ -16326,7 +16326,7 @@ static ggml_type llama_tensor_get_type(quantize_state_internal & qs, ggml_type n
// same quantization as before imatrix stuff, and b) Q4_1/Q5_1 do go crazy on ffn_down without an imatrix.
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) {
else if (ftype == LLAMA_FTYPE_MOSTLY_Q4_0_R8 && qs.has_imatrix && i_layer < n_layer/8) {
new_type = GGML_TYPE_IQ4_NL_R4;
}
++qs.i_ffn_down;
@ -16339,7 +16339,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_Q4_K_R4 ||
ftype == LLAMA_FTYPE_MOSTLY_IQ4_NL_R4 || ftype == LLAMA_FTYPE_MOSTLY_IQ4_XS_R4 || ftype == LLAMA_FTYPE_MOSTLY_Q3_K_R4 ||
ftype == LLAMA_FTYPE_MOSTLY_IQ4_NL_R4 || ftype == LLAMA_FTYPE_MOSTLY_IQ4_XS_R8 || ftype == LLAMA_FTYPE_MOSTLY_Q3_K_R4 ||
ftype == LLAMA_FTYPE_MOSTLY_Q2_K_R4|| ftype == LLAMA_FTYPE_MOSTLY_IQ4_K_R4 || ftype == LLAMA_FTYPE_MOSTLY_IQ3_K_R4 ||
ftype == LLAMA_FTYPE_MOSTLY_IQ2_K_R4 || ftype == LLAMA_FTYPE_MOSTLY_IQ3_XXS_R4 || ftype == LLAMA_FTYPE_MOSTLY_IQ3_S_R4) {
new_type = GGML_TYPE_Q5_K;
@ -16411,7 +16411,7 @@ static ggml_type llama_tensor_get_type(quantize_state_internal & qs, ggml_type n
new_type == GGML_TYPE_IQ3_XXS || new_type == GGML_TYPE_IQ1_S || new_type == GGML_TYPE_IQ3_S ||
new_type == GGML_TYPE_IQ1_M || new_type == GGML_TYPE_IQ4_K || new_type == GGML_TYPE_IQ2_K ||
new_type == GGML_TYPE_IQ5_K || new_type == GGML_TYPE_IQ3_K || new_type == GGML_TYPE_Q4_K_R4 ||
new_type == GGML_TYPE_IQ6_K || new_type == GGML_TYPE_IQ4_KS || new_type == GGML_TYPE_IQ4_XS_R4 ||
new_type == GGML_TYPE_IQ6_K || new_type == GGML_TYPE_IQ4_KS || new_type == GGML_TYPE_IQ4_XS_R8 ||
new_type == GGML_TYPE_IQ2_KS || new_type == GGML_TYPE_IQ4_KSS || new_type == GGML_TYPE_Q6_K_R4 ||
new_type == GGML_TYPE_Q5_K_R4 || new_type == GGML_TYPE_Q3_K_R4 || new_type == GGML_TYPE_Q2_K_R4 ||
new_type == GGML_TYPE_IQ4_K_R4|| new_type == GGML_TYPE_Q8_K_R8 || new_type == GGML_TYPE_IQ3_K_R4||
@ -16459,7 +16459,7 @@ static ggml_type llama_tensor_get_type(quantize_state_internal & qs, ggml_type n
case GGML_TYPE_IQ4_KSS:
case GGML_TYPE_IQ4_KS:
case GGML_TYPE_IQ4_KS_R4:
case GGML_TYPE_IQ4_XS_R4:
case GGML_TYPE_IQ4_XS_R8:
case GGML_TYPE_IQ4_XS: new_type = GGML_TYPE_IQ4_NL; break;
case GGML_TYPE_IQ4_K:
case GGML_TYPE_IQ4_K_R4:
@ -16589,11 +16589,11 @@ static void llama_model_quantize_internal(const std::string & fname_inp, const s
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_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_IQ4_XS_R8:default_type = GGML_TYPE_IQ4_XS_R8;break;
case LLAMA_FTYPE_MOSTLY_Q4_0_R8: default_type = GGML_TYPE_Q4_0_R8; break;
case LLAMA_FTYPE_MOSTLY_Q5_0_R4: default_type = GGML_TYPE_Q5_0_R4; break;
case LLAMA_FTYPE_MOSTLY_Q6_0_R4: default_type = GGML_TYPE_Q6_0_R4; break;
case LLAMA_FTYPE_MOSTLY_Q8_0_R4: default_type = GGML_TYPE_Q8_0_R4; break;
case LLAMA_FTYPE_MOSTLY_Q8_0_R8: default_type = GGML_TYPE_Q8_0_R8; break;
case LLAMA_FTYPE_MOSTLY_IQ4_XS: default_type = GGML_TYPE_IQ4_XS; break;
case LLAMA_FTYPE_MOSTLY_IQ4_KS: default_type = GGML_TYPE_IQ4_KS; break;
case LLAMA_FTYPE_MOSTLY_IQ4_KS_R4:default_type = GGML_TYPE_IQ4_KS_R4;break;
@ -16969,11 +16969,11 @@ static void llama_model_quantize_internal(const std::string & fname_inp, const s
if (tensor->ne[1] % 4 != 0) new_type = GGML_TYPE_IQ4_NL;
else chunk_size_multiplier = 4;
}
else if (new_type == GGML_TYPE_IQ4_XS_R4) {
else if (new_type == GGML_TYPE_IQ4_XS_R8) {
if (tensor->ne[1] % 8 != 0) new_type = GGML_TYPE_IQ4_XS;
else chunk_size_multiplier = 8;
}
else if (new_type == GGML_TYPE_Q4_0_R4) {
else if (new_type == GGML_TYPE_Q4_0_R8) {
if (tensor->ne[1] % 8 != 0) new_type = GGML_TYPE_Q4_0;
else chunk_size_multiplier = 8;
}
@ -16985,7 +16985,7 @@ static void llama_model_quantize_internal(const std::string & fname_inp, const s
if (tensor->ne[1] % 4 != 0) new_type = GGML_TYPE_Q6_0;
else chunk_size_multiplier = 4;
}
else if (new_type == GGML_TYPE_Q8_0_R4) {
else if (new_type == GGML_TYPE_Q8_0_R8) {
if (tensor->ne[1] % 8 != 0) new_type = GGML_TYPE_Q8_0;
else chunk_size_multiplier = 8;
}