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

Convert models to row-interleaved quants using the quantize tool (#272) 

* Repack a model with the quantize tool

* WIP

* Fixed various issues

As we don't have a way to tell if a repacked quant has been modified,
I had to remove the modification at the expense of a slight decrease
in performance. This affects q8_0_r8, q8_KV_r8, q8_k_r8 on Zen4, and
q4_0_r8 on ARM.

* Create wk_b and wv_b as Q8_0_R8 if the wkv_b type is interleaved

* Fix GCC 13.3 compilation error

* Another one

* Add missing include

---------

Co-authored-by: Iwan Kawrakow <iwan.kawrakow@gmail.com>
This commit is contained in:
Kawrakow 2025-03-21 07:23:36 +01:00 committed by GitHub
parent 712de34b12
commit c5e554f941
7 changed files with 246 additions and 58 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
examples/quantize/quantize.cpp
View File

@ -145,6 +145,7 @@ static void usage(const char * executable) {
printf(" --output-tensor-type ggml_type: use this ggml_type for the output.weight tensor.\n"); printf(" --output-tensor-type ggml_type: use this ggml_type for the output.weight tensor.\n");
printf(" --token-embedding-type ggml_type: use this ggml_type for the token_embd.weight tensor.\n\n"); printf(" --token-embedding-type ggml_type: use this ggml_type for the token_embd.weight tensor.\n\n");
printf(" --custom-q regex1=type1,regex2=type2...: use this to specify custom quantization type rules.\n\n"); printf(" --custom-q regex1=type1,regex2=type2...: use this to specify custom quantization type rules.\n\n");
printf(" --repack Repack all tensors to the corresponding _r4/8 variant if available.\n\n");
printf("Additional specific tensor quantization types used in the custom quant scheme 'CQS (default is Q2_K):\n"); printf("Additional specific tensor quantization types used in the custom quant scheme 'CQS (default is Q2_K):\n");
printf(" --attn-q-type ggml_type: use this ggml_type for the attn_q.weight tensor.\n"); printf(" --attn-q-type ggml_type: use this ggml_type for the attn_q.weight tensor.\n");
printf(" --attn-k-type ggml_type: use this ggml_type for the attn_k.weight tensor.\n"); printf(" --attn-k-type ggml_type: use this ggml_type for the attn_k.weight tensor.\n");
@ -331,6 +332,8 @@ int main(int argc, char ** argv) {
params.quantize_output_tensor = false; params.quantize_output_tensor = false;
} else if (strcmp(argv[arg_idx], "--ignore-imatrix-rules") == 0) { } else if (strcmp(argv[arg_idx], "--ignore-imatrix-rules") == 0) {
params.ignore_imatrix_rules = true; params.ignore_imatrix_rules = true;
} else if (strcmp(argv[arg_idx], "--repack") == 0) {
params.only_repack = true;
} else if (strcmp(argv[arg_idx], "--output-tensor-type") == 0) { } else if (strcmp(argv[arg_idx], "--output-tensor-type") == 0) {
if (arg_idx < argc-1) { if (arg_idx < argc-1) {
params.output_tensor_type = parse_ggml_type(argv[++arg_idx]); params.output_tensor_type = parse_ggml_type(argv[++arg_idx]);

48
ggml/src/ggml.c
View File

@ -1756,6 +1756,15 @@ ggml_type_traits_t ggml_internal_get_type_traits(enum ggml_type type) {
return type_traits[type]; return type_traits[type];
} }
static inline int ggml_packed_rows(enum ggml_type type) {
return type == GGML_TYPE_BF16_R16 ? 16
: type == GGML_TYPE_Q8_K_R8 || type == GGML_TYPE_Q8_KV_R8 ||
type == GGML_TYPE_Q8_0_R8 || type == GGML_TYPE_Q4_0_R8 ||
type == GGML_TYPE_IQ4_XS_R8 ? 8
: type >= GGML_TYPE_Q4_0_R8 && type <= GGML_TYPE_Q8_K_R8 ? 4
: 1;
}
// //
// simd mappings // simd mappings
// //
@ -10119,9 +10128,11 @@ static void ggml_compute_forward_dup_f32(
} }
// parallelize by rows // parallelize by rows
int n_packed = ggml_packed_rows(dst->type);
GGML_ASSERT(dst->ne[1] % n_packed == 0);
const int nr = ne01; const int nr = ne01;
// number of rows per thread // number of rows per thread
const int dr = (nr + nth - 1) / nth; const int dr = n_packed*((nr/n_packed + nth - 1) / nth);
// row range for this thread // row range for this thread
const int ir0 = dr * ith; const int ir0 = dr * ith;
const int ir1 = MIN(ir0 + dr, nr); const int ir1 = MIN(ir0 + dr, nr);
@ -10173,10 +10184,10 @@ static void ggml_compute_forward_dup_f32(
for (int i03 = 0; i03 < ne03; i03++) { for (int i03 = 0; i03 < ne03; i03++) {
for (int i02 = 0; i02 < ne02; i02++) { for (int i02 = 0; i02 < ne02; i02++) {
id += rs * ir0; id += rs * ir0;
for (int i01 = ir0; i01 < ir1; i01++) { for (int i01 = ir0; i01 < ir1; i01 += n_packed) {
const float * src0_ptr = (float *) ((char *) src0->data + i01*nb01 + i02*nb02 + i03*nb03); const float * src0_ptr = (float *) ((char *) src0->data + i01*nb01 + i02*nb02 + i03*nb03);
quantize_row_q(src0_ptr, dst_ptr + id, ne00); quantize_row_q(src0_ptr, dst_ptr + id, ne00*n_packed);
id += rs; id += rs*n_packed;
} }
id += rs * (ne01 - ir1); id += rs * (ne01 - ir1);
} }
@ -10441,10 +10452,15 @@ static void ggml_compute_forward_dup_bytes(
// parallelize by rows // parallelize by rows
const int nr = ne01; const int nr = ne01;
const int n_packed = ggml_packed_rows(dst->type);
GGML_ASSERT(nr%n_packed == 0);
const int nrp = nr/n_packed;
// number of rows per thread // number of rows per thread
const int dr = (nr + nth - 1) / nth; const int drp = (nrp + nth - 1) / nth;
const int dr = drp*n_packed;
// row range for this thread // row range for this thread
const int ir0 = dr * ith; const int ir0 = dr * ith;
if (ir0 >= nr) return;
const int ir1 = MIN(ir0 + dr, nr); const int ir1 = MIN(ir0 + dr, nr);
if (src0->type == dst->type && if (src0->type == dst->type &&
@ -10569,10 +10585,19 @@ static void ggml_compute_forward_dup_q(
const struct ggml_compute_params * params, const struct ggml_compute_params * params,
struct ggml_tensor * dst) { struct ggml_tensor * dst) {
GGML_ASSERT(ggml_is_quantized(dst->src[0]->type));
int64_t nrows = ggml_nrows(dst); int64_t nrows = ggml_nrows(dst);
int ith = params->ith; int ith = params->ith;
int nth = params->nth; int nth = params->nth;
if (dst->src[0]->type == dst->type &&
dst->src[0]->nb[0] == ggml_type_size(dst->type) &&
dst->nb[0] == ggml_type_size(dst->type)) {
ggml_compute_forward_dup_bytes(params, dst);
return;
}
if (dst->type == GGML_TYPE_Q8_0 && dst->src[0]->type == GGML_TYPE_Q8_0 && if (dst->type == GGML_TYPE_Q8_0 && dst->src[0]->type == GGML_TYPE_Q8_0 &&
ggml_are_same_shape(dst, dst->src[0])) { ggml_are_same_shape(dst, dst->src[0])) {
@ -10626,6 +10651,10 @@ static void ggml_compute_forward_dup_q(
return; return;
} }
if (dst->type != GGML_TYPE_F32) {
printf("%s: %s -> %s is of type %s\n", __func__, dst->src[0]->name, dst->name, ggml_type_name(dst->type));
GGML_ABORT("fatal error");
}
GGML_ASSERT(dst->type == GGML_TYPE_F32); GGML_ASSERT(dst->type == GGML_TYPE_F32);
struct ggml_tensor * src0 = dst->src[0]; struct ggml_tensor * src0 = dst->src[0];
GGML_ASSERT(src0->ne[0] == dst->ne[0] && src0->nb[0] == ggml_type_size(src0->type)); GGML_ASSERT(src0->ne[0] == dst->ne[0] && src0->nb[0] == ggml_type_size(src0->type));
@ -10633,12 +10662,15 @@ static void ggml_compute_forward_dup_q(
ggml_to_float_t to_float = type_traits[src0->type].to_float; ggml_to_float_t to_float = type_traits[src0->type].to_float;
GGML_ASSERT(to_float != NULL); GGML_ASSERT(to_float != NULL);
int64_t n_per_thread = (nrows + nth - 1)/nth; int n_packed = ggml_packed_rows(src0->type);
GGML_ASSERT(src0->ne[1] % n_packed == 0);
int64_t n_per_thread = n_packed*((nrows/n_packed + nth - 1)/nth);
int64_t first_row = ith*n_per_thread; int64_t first_row = ith*n_per_thread;
if (first_row >= nrows) return; if (first_row >= nrows) return;
int64_t last_row = MIN(first_row + n_per_thread, nrows); int64_t last_row = MIN(first_row + n_per_thread, nrows);
for (int64_t ir = first_row; ir < last_row; ++ir) { for (int64_t ir = first_row; ir < last_row; ir += n_packed) {
int64_t i03 = ir/(src0->ne[1]*src0->ne[2]); int64_t i03 = ir/(src0->ne[1]*src0->ne[2]);
int64_t i02 = (ir - i03*src0->ne[1]*src0->ne[2])/src0->ne[1]; int64_t i02 = (ir - i03*src0->ne[1]*src0->ne[2])/src0->ne[1];
int64_t i01 = ir - i03*src0->ne[1]*src0->ne[2] - i02*src0->ne[1]; int64_t i01 = ir - i03*src0->ne[1]*src0->ne[2] - i02*src0->ne[1];
@ -10649,7 +10681,7 @@ static void ggml_compute_forward_dup_q(
const char * q = (const char *)src0->data + i03*src0->nb[3] + i02*src0->nb[2] + i01*src0->nb[1]; const char * q = (const char *)src0->data + i03*src0->nb[3] + i02*src0->nb[2] + i01*src0->nb[1];
char * f = ( char *)dst->data + i3* dst->nb[3] + i2* dst->nb[2] + i1* dst->nb[1]; char * f = ( char *)dst->data + i3* dst->nb[3] + i2* dst->nb[2] + i1* dst->nb[1];
to_float((const void *)q, (float *)f, src0->ne[0]); to_float((const void *)q, (float *)f, src0->ne[0]*n_packed);
} }
} }

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

@ -4434,14 +4434,17 @@ inline __m256i qx_r8_q8_dot_product(const __m256i * qx, const int8_t * y) {
return sumi; return sumi;
} }
inline __m256i q8_0_r8_dot_product(const uint8_t * x, const int8_t * y, __m256i * qx) { inline __m256i q8_0_r8_dot_product(const uint8_t * x, const int8_t * y, __m256i * qx) {
qx[0] = _mm256_loadu_si256((const __m256i *)x+0); for (int i = 0; i < 8; ++i) {
qx[1] = _mm256_loadu_si256((const __m256i *)x+1); qx[i] = _mm256_add_epi8(_mm256_loadu_si256((const __m256i *)x+i), _mm256_set1_epi8(127));
qx[2] = _mm256_loadu_si256((const __m256i *)x+2); }
qx[3] = _mm256_loadu_si256((const __m256i *)x+3); //qx[0] = _mm256_loadu_si256((const __m256i *)x+0);
qx[4] = _mm256_loadu_si256((const __m256i *)x+4); //qx[1] = _mm256_loadu_si256((const __m256i *)x+1);
qx[5] = _mm256_loadu_si256((const __m256i *)x+5); //qx[2] = _mm256_loadu_si256((const __m256i *)x+2);
qx[6] = _mm256_loadu_si256((const __m256i *)x+6); //qx[3] = _mm256_loadu_si256((const __m256i *)x+3);
qx[7] = _mm256_loadu_si256((const __m256i *)x+7); //qx[4] = _mm256_loadu_si256((const __m256i *)x+4);
//qx[5] = _mm256_loadu_si256((const __m256i *)x+5);
//qx[6] = _mm256_loadu_si256((const __m256i *)x+6);
//qx[7] = _mm256_loadu_si256((const __m256i *)x+7);
return qx_r8_q8_dot_product(qx, y); return qx_r8_q8_dot_product(qx, y);
} }
template <int nrc_y> template <int nrc_y>
@ -4496,6 +4499,7 @@ static void mul_mat_q8_0_r8_q8_1(int n, const void * vx, size_t bx, const DataIn
for (int j = 0; j < 8; ++j) { for (int j = 0; j < 8; ++j) {
qx[j] = _mm512_inserti32x8(_mm512_castsi256_si512(_mm256_loadu_si256((const __m256i *)q8l[4*ib4+k].qs+j)), qx[j] = _mm512_inserti32x8(_mm512_castsi256_si512(_mm256_loadu_si256((const __m256i *)q8l[4*ib4+k].qs+j)),
_mm256_loadu_si256((const __m256i *)q8h[4*ib4+k].qs+j), 1); _mm256_loadu_si256((const __m256i *)q8h[4*ib4+k].qs+j), 1);
qx[j] = _mm512_add_epi8(qx[j], _mm512_set1_epi8(127));
} }
for (int iy = 0; iy < nrc_y; ++iy) { for (int iy = 0; iy < nrc_y; ++iy) {
auto sumi = qx_r8_q8_dot_product(qx, q8.y[iy][ib4].qs+32*k); auto sumi = qx_r8_q8_dot_product(qx, q8.y[iy][ib4].qs+32*k);
@ -4512,6 +4516,7 @@ static void mul_mat_q8_0_r8_q8_1(int n, const void * vx, size_t bx, const DataIn
for (int j = 0; j < 8; ++j) { for (int j = 0; j < 8; ++j) {
qx[j] = _mm512_inserti32x8(_mm512_castsi256_si512(_mm256_loadu_si256((const __m256i *)q8l[ib].qs+j)), qx[j] = _mm512_inserti32x8(_mm512_castsi256_si512(_mm256_loadu_si256((const __m256i *)q8l[ib].qs+j)),
_mm256_loadu_si256((const __m256i *)q8h[ib].qs+j), 1); _mm256_loadu_si256((const __m256i *)q8h[ib].qs+j), 1);
qx[j] = _mm512_add_epi8(qx[j], _mm512_set1_epi8(127));
} }
for (int iy = 0; iy < nrc_y; ++iy) { for (int iy = 0; iy < nrc_y; ++iy) {
auto qy = (const block_q8_1 *)q8.y[iy]; auto qy = (const block_q8_1 *)q8.y[iy];
@ -6347,6 +6352,11 @@ static void mul_mat_q8_k_r8_q8_k(int n, const void * vx, size_t bx, const DataIn
auto s1 = _mm256_sign_epi8(qx[1], qx[1]); auto s1 = _mm256_sign_epi8(qx[1], qx[1]);
auto s2 = _mm256_sign_epi8(qx[2], qx[2]); auto s2 = _mm256_sign_epi8(qx[2], qx[2]);
auto s3 = _mm256_sign_epi8(qx[3], qx[3]); auto s3 = _mm256_sign_epi8(qx[3], qx[3]);
#else
qx[0] = _mm256_add_epi8(qx[0], _mm256_set1_epi8(127));
qx[1] = _mm256_add_epi8(qx[1], _mm256_set1_epi8(127));
qx[2] = _mm256_add_epi8(qx[2], _mm256_set1_epi8(127));
qx[3] = _mm256_add_epi8(qx[3], _mm256_set1_epi8(127));
#endif #endif
for (int iy = 0; iy < nrc_y; ++iy) { for (int iy = 0; iy < nrc_y; ++iy) {
auto y128 = _mm_loadu_si128((const __m128i*)q8.y[iy][ibl].qs+ib); auto y128 = _mm_loadu_si128((const __m128i*)q8.y[iy][ibl].qs+ib);
@ -6425,6 +6435,11 @@ static void mul_mat_q8_KV_r8_q8_KV(int n, const void * vx, size_t bx, const Data
auto s1 = _mm256_sign_epi8(qx[1], qx[1]); auto s1 = _mm256_sign_epi8(qx[1], qx[1]);
auto s2 = _mm256_sign_epi8(qx[2], qx[2]); auto s2 = _mm256_sign_epi8(qx[2], qx[2]);
auto s3 = _mm256_sign_epi8(qx[3], qx[3]); auto s3 = _mm256_sign_epi8(qx[3], qx[3]);
#else
qx[0] = _mm256_add_epi8(qx[0], _mm256_set1_epi8(127));
qx[1] = _mm256_add_epi8(qx[1], _mm256_set1_epi8(127));
qx[2] = _mm256_add_epi8(qx[2], _mm256_set1_epi8(127));
qx[3] = _mm256_add_epi8(qx[3], _mm256_set1_epi8(127));
#endif #endif
for (int iy = 0; iy < nrc_y; ++iy) { for (int iy = 0; iy < nrc_y; ++iy) {
auto y128 = _mm_loadu_si128((const __m128i*)q8y[iy]+ib); auto y128 = _mm_loadu_si128((const __m128i*)q8y[iy]+ib);
@ -14305,8 +14320,8 @@ struct Q4_0_R8_Dequantizer {
float32x4x2_t scales = { vcvt_f32_f16(vget_low_f16(scales16)), vcvt_f32_f16(vget_high_f16(scales16)) }; float32x4x2_t scales = { vcvt_f32_f16(vget_low_f16(scales16)), vcvt_f32_f16(vget_high_f16(scales16)) };
for (int j = 0; j < 4; ++j) { for (int j = 0; j < 4; ++j) {
auto bits = vld1q_u8_x2(iq4[4*ib4+k].qs + 32*j); auto bits = vld1q_u8_x2(iq4[4*ib4+k].qs + 32*j);
//bits.val[0] = veorq_u8(m88, bits.val[0]); bits.val[0] = veorq_u8(m88, bits.val[0]);
//bits.val[1] = veorq_u8(m88, bits.val[1]); bits.val[1] = veorq_u8(m88, bits.val[1]);
qx[2*j+0] = vshlq_n_u8(bits.val[0], 4); qx[2*j+0] = vshlq_n_u8(bits.val[0], 4);
qx[2*j+1] = vandq_u8(bits.val[0], m4); qx[2*j+1] = vandq_u8(bits.val[0], m4);
qx[2*j+8] = vshlq_n_u8(bits.val[1], 4); qx[2*j+8] = vshlq_n_u8(bits.val[1], 4);
@ -15305,12 +15320,12 @@ struct HelperQ80R8 : public BaseHelper<step> {
m1 = _mm256_unpackhi_epi64(t0, t1); m1 = _mm256_unpackhi_epi64(t0, t1);
m2 = _mm256_unpacklo_epi64(t2, t3); m2 = _mm256_unpacklo_epi64(t2, t3);
m3 = _mm256_unpackhi_epi64(t2, t3); m3 = _mm256_unpackhi_epi64(t2, t3);
#ifdef HAVE_FANCY_SIMD //#ifdef HAVE_FANCY_SIMD
m0 = _mm256_add_epi8(m0, _mm256_set1_epi8(127)); // m0 = _mm256_add_epi8(m0, _mm256_set1_epi8(127));
m1 = _mm256_add_epi8(m1, _mm256_set1_epi8(127)); // m1 = _mm256_add_epi8(m1, _mm256_set1_epi8(127));
m2 = _mm256_add_epi8(m2, _mm256_set1_epi8(127)); // m2 = _mm256_add_epi8(m2, _mm256_set1_epi8(127));
m3 = _mm256_add_epi8(m3, _mm256_set1_epi8(127)); // m3 = _mm256_add_epi8(m3, _mm256_set1_epi8(127));
#endif //#endif
_mm256_storeu_si256((__m256i *)y[ib].qs + 0, m0); _mm256_storeu_si256((__m256i *)y[ib].qs + 0, m0);
_mm256_storeu_si256((__m256i *)y[ib].qs + 1, m1); _mm256_storeu_si256((__m256i *)y[ib].qs + 1, m1);
_mm256_storeu_si256((__m256i *)y[ib].qs + 2, m2); _mm256_storeu_si256((__m256i *)y[ib].qs + 2, m2);
@ -15327,12 +15342,12 @@ struct HelperQ80R8 : public BaseHelper<step> {
m1 = _mm256_unpackhi_epi64(t0, t1); m1 = _mm256_unpackhi_epi64(t0, t1);
m2 = _mm256_unpacklo_epi64(t2, t3); m2 = _mm256_unpacklo_epi64(t2, t3);
m3 = _mm256_unpackhi_epi64(t2, t3); m3 = _mm256_unpackhi_epi64(t2, t3);
#ifdef HAVE_FANCY_SIMD //#ifdef HAVE_FANCY_SIMD
m0 = _mm256_add_epi8(m0, _mm256_set1_epi8(127)); // m0 = _mm256_add_epi8(m0, _mm256_set1_epi8(127));
m1 = _mm256_add_epi8(m1, _mm256_set1_epi8(127)); // m1 = _mm256_add_epi8(m1, _mm256_set1_epi8(127));
m2 = _mm256_add_epi8(m2, _mm256_set1_epi8(127)); // m2 = _mm256_add_epi8(m2, _mm256_set1_epi8(127));
m3 = _mm256_add_epi8(m3, _mm256_set1_epi8(127)); // m3 = _mm256_add_epi8(m3, _mm256_set1_epi8(127));
#endif //#endif
_mm256_storeu_si256((__m256i *)y[ib].qs + 4, m0); _mm256_storeu_si256((__m256i *)y[ib].qs + 4, m0);
_mm256_storeu_si256((__m256i *)y[ib].qs + 5, m1); _mm256_storeu_si256((__m256i *)y[ib].qs + 5, m1);
_mm256_storeu_si256((__m256i *)y[ib].qs + 6, m2); _mm256_storeu_si256((__m256i *)y[ib].qs + 6, m2);
@ -15424,12 +15439,12 @@ struct HelperQ8KVR8 : public BaseHelper<step> {
m1 = _mm256_unpackhi_epi64(t0, t1); m1 = _mm256_unpackhi_epi64(t0, t1);
m2 = _mm256_unpacklo_epi64(t2, t3); m2 = _mm256_unpacklo_epi64(t2, t3);
m3 = _mm256_unpackhi_epi64(t2, t3); m3 = _mm256_unpackhi_epi64(t2, t3);
#ifdef HAVE_FANCY_SIMD //#ifdef HAVE_FANCY_SIMD
m0 = _mm256_add_epi8(m0, _mm256_set1_epi8(127)); // m0 = _mm256_add_epi8(m0, _mm256_set1_epi8(127));
m1 = _mm256_add_epi8(m1, _mm256_set1_epi8(127)); // m1 = _mm256_add_epi8(m1, _mm256_set1_epi8(127));
m2 = _mm256_add_epi8(m2, _mm256_set1_epi8(127)); // m2 = _mm256_add_epi8(m2, _mm256_set1_epi8(127));
m3 = _mm256_add_epi8(m3, _mm256_set1_epi8(127)); // m3 = _mm256_add_epi8(m3, _mm256_set1_epi8(127));
#endif //#endif
_mm256_storeu_si256((__m256i *)y[ix].qs + 4*ib+0, m0); _mm256_storeu_si256((__m256i *)y[ix].qs + 4*ib+0, m0);
_mm256_storeu_si256((__m256i *)y[ix].qs + 4*ib+1, m1); _mm256_storeu_si256((__m256i *)y[ix].qs + 4*ib+1, m1);
_mm256_storeu_si256((__m256i *)y[ix].qs + 4*ib+2, m2); _mm256_storeu_si256((__m256i *)y[ix].qs + 4*ib+2, m2);

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

@ -25,6 +25,7 @@
#include <thread> #include <thread>
#include <atomic> #include <atomic>
#include <unordered_map> #include <unordered_map>
#include <string>
#if defined(_MSC_VER) #if defined(_MSC_VER)
#pragma warning(disable: 4244 4267) // possible loss of data #pragma warning(disable: 4244 4267) // possible loss of data
@ -6766,9 +6767,7 @@ struct Modify {
modify_func_t mod_func; modify_func_t mod_func;
int nrows; int nrows;
}; };
} const Modify * get_modify_info(ggml_type type) {
bool iqk_modify_tensor(struct ggml_tensor * tensor) {
static const std::unordered_map<ggml_type, Modify> k_mod_map = { static const std::unordered_map<ggml_type, Modify> k_mod_map = {
#ifdef __ARM_NEON #ifdef __ARM_NEON
{ GGML_TYPE_Q4_0_R8, {modify_q4_0_r8, 8} }, { GGML_TYPE_Q4_0_R8, {modify_q4_0_r8, 8} },
@ -6779,10 +6778,31 @@ bool iqk_modify_tensor(struct ggml_tensor * tensor) {
{ GGML_TYPE_Q8_KV_R8, {modify_q8_KV_r8, 8} }, { GGML_TYPE_Q8_KV_R8, {modify_q8_KV_r8, 8} },
#endif #endif
}; };
auto it = k_mod_map.find(tensor->type); auto it = k_mod_map.find(type);
if (it == k_mod_map.end()) return false; return it != k_mod_map.end() ? &it->second : nullptr;
}
bool is_forbidden_tensor(const std::string& name) {
static const std::string kTokenEmbd{"token_embd.weight"};
if (name == kTokenEmbd) return true;
//if (auto pos = name.find("attn_kv_b.weight"); pos != std::string::npos) return true;
return false;
}
}
auto& m = it->second; bool iqk_should_modify_tensor([[maybe_unused]] const struct ggml_tensor * tensor) {
return false;
//if (is_forbidden_tensor(tensor->name)) return false;
//auto mptr = get_modify_info(tensor->type);
//return mptr ? true : false;
}
bool iqk_modify_tensor(struct ggml_tensor * tensor) {
return false;
auto mptr = get_modify_info(tensor->type);
if (!mptr) return false;
if (is_forbidden_tensor(std::string{tensor->name})) return false;
auto& m = *mptr;
int nrows = ggml_nrows(tensor); int nrows = ggml_nrows(tensor);
int nchunks = nrows/m.nrows; int nchunks = nrows/m.nrows;
int max_thread = std::max(1, int(std::thread::hardware_concurrency()/2)); int max_thread = std::max(1, int(std::thread::hardware_concurrency()/2));
@ -6805,12 +6825,8 @@ bool iqk_modify_tensor(struct ggml_tensor * tensor) {
return true; return true;
} }
void iqk_repack_tensor(struct ggml_tensor * tensor) { namespace {
constexpr int kChunk = 8; const Repack * get_repack_info(ggml_type type) {
if (!tensor) return;
if (!ggml_is_contiguous(tensor)) return;
if (strncmp(tensor->name, "token_embd.weight", GGML_MAX_NAME) == 0) return;
if (tensor->ne[1] % 4) return;
static const std::unordered_map<ggml_type, Repack> k_map = { static const std::unordered_map<ggml_type, Repack> k_map = {
{ GGML_TYPE_IQ2_K, { GGML_TYPE_IQ2_K_R4, 4, (Repack::repack_func)repack_iq2_k} }, { GGML_TYPE_IQ2_K, { GGML_TYPE_IQ2_K_R4, 4, (Repack::repack_func)repack_iq2_k} },
{ GGML_TYPE_IQ3_K, { GGML_TYPE_IQ3_K_R4, 4, (Repack::repack_func)repack_iq3_k} }, { GGML_TYPE_IQ3_K, { GGML_TYPE_IQ3_K_R4, 4, (Repack::repack_func)repack_iq3_k} },
@ -6841,12 +6857,30 @@ void iqk_repack_tensor(struct ggml_tensor * tensor) {
{ GGML_TYPE_F16, { GGML_TYPE_BF16_R16, 16, (Repack::repack_func)repack_bf16<ggml_half>} }, { GGML_TYPE_F16, { GGML_TYPE_BF16_R16, 16, (Repack::repack_func)repack_bf16<ggml_half>} },
#endif #endif
}; };
auto it = k_map.find(type);
return it != k_map.end() ? &it->second : nullptr;
}
}
auto it = k_map.find(tensor->type); int iqk_repacked_type(const struct ggml_tensor * tensor) {
if (it == k_map.end()) return; if (!ggml_is_contiguous(tensor)) return (int)tensor->type;
if (tensor->ne[1] % it->second.num_rows) return; if (is_forbidden_tensor(tensor->name)) return (int)tensor->type;
auto rptr = get_repack_info(tensor->type);
return rptr && tensor->ne[1] % rptr->num_rows == 0 ? (int)rptr->new_type : (int)tensor->type;
}
auto& r = it->second; void iqk_repack_tensor(struct ggml_tensor * tensor) {
constexpr int kChunk = 8;
if (!tensor) return;
if (!ggml_is_contiguous(tensor)) return;
if (is_forbidden_tensor(tensor->name)) return;
if (tensor->ne[1] % 4) return;
auto rptr = get_repack_info(tensor->type);
if (!rptr) return;
if (tensor->ne[1] % rptr->num_rows) return;
auto& r = *rptr;
auto nrows = ggml_nrows(tensor); auto nrows = ggml_nrows(tensor);
@ -6871,7 +6905,8 @@ void iqk_repack_tensor(struct ggml_tensor * tensor) {
int last_row = std::min(first_row + chunkSize*r.num_rows, nrows); int last_row = std::min(first_row + chunkSize*r.num_rows, nrows);
for (int row = first_row; row < last_row; row += r.num_rows) { for (int row = first_row; row < last_row; row += r.num_rows) {
std::memcpy(qtmp.data(), data + row*row_size, r.num_rows*row_size); std::memcpy(qtmp.data(), data + row*row_size, r.num_rows*row_size);
r.repack(r.num_rows, n_per_row, qtmp.data(), data + row*row_size, true); //r.repack(r.num_rows, n_per_row, qtmp.data(), data + row*row_size, true);
r.repack(r.num_rows, n_per_row, qtmp.data(), data + row*row_size, false);
} }
} }
}; };

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

@ -245,6 +245,9 @@ void repack_bf16_bf16_r16(const void * GGML_RESTRICT src, void * GGML_RESTRICT d
void iqk_repack_tensor(struct ggml_tensor * tensor); void iqk_repack_tensor(struct ggml_tensor * tensor);
bool iqk_modify_tensor(struct ggml_tensor * tensor); bool iqk_modify_tensor(struct ggml_tensor * tensor);
int iqk_repacked_type(const struct ggml_tensor * tensor); // int instead of ggml_type so we don't need to include ggml.h
bool iqk_should_modify_tensor(const struct ggml_tensor * tensor);
// So we can re-pack Microsoft's BitNet I2_S quants // So we can re-pack Microsoft's BitNet I2_S quants
void dequantize_row_ms_i2s(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); void dequantize_row_ms_i2s(const void * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);

1
include/llama.h
View File

@ -416,6 +416,7 @@ extern "C" {
bool pure; // quantize all tensors to the default type bool pure; // quantize all tensors to the default type
bool keep_split; // quantize to the same number of shards bool keep_split; // quantize to the same number of shards
bool ignore_imatrix_rules; // If set to true, the built-in rules for refusing to quantize into certain quants without imatrix are ignored bool ignore_imatrix_rules; // If set to true, the built-in rules for refusing to quantize into certain quants without imatrix are ignored
bool only_repack; // Only repack tensors
void * imatrix; // pointer to importance matrix data void * imatrix; // pointer to importance matrix data
void * kv_overrides; // pointer to vector containing overrides void * kv_overrides; // pointer to vector containing overrides
void * custom_quants; // pointer to vector containing custom quantization rules void * custom_quants; // pointer to vector containing custom quantization rules

109
src/llama.cpp
View File

@ -8194,7 +8194,8 @@ static bool llm_load_tensors(
auto wk_b_f32_t = ggml_cont(ctx, wk_b_f32_tview); auto wk_b_f32_t = ggml_cont(ctx, wk_b_f32_tview);
wk_b_f32_t->data = (char *)wk_b_f32->data + ggml_nbytes(wk_b_f32); wk_b_f32_t->data = (char *)wk_b_f32->data + ggml_nbytes(wk_b_f32);
auto new_type = ggml_is_quantized(wkv_b.type) ? GGML_TYPE_Q8_0 : wkv_b.type; auto new_type = ggml_is_quantized(wkv_b.type) ?
wkv_b.type >= GGML_TYPE_Q4_0_R8 && wkv_b.type <= GGML_TYPE_Q8_K_R8 ? GGML_TYPE_Q8_0_R8 : GGML_TYPE_Q8_0 : wkv_b.type;
auto wk_b = ggml_cast(ctx, wk_b_f32_t, new_type); auto wk_b = ggml_cast(ctx, wk_b_f32_t, new_type);
wk_b->data = (char *)wk_b_f32_t->data + ggml_nbytes(wk_b_f32_t); wk_b->data = (char *)wk_b_f32_t->data + ggml_nbytes(wk_b_f32_t);
@ -8218,6 +8219,9 @@ static bool llm_load_tensors(
ggml_set_name(l.computed_wk_b.get(), name.c_str()); ggml_set_name(l.computed_wk_b.get(), name.c_str());
ggml_backend_buffer_set_usage(l.computed_wk_b->buffer, GGML_BACKEND_BUFFER_USAGE_WEIGHTS); ggml_backend_buffer_set_usage(l.computed_wk_b->buffer, GGML_BACKEND_BUFFER_USAGE_WEIGHTS);
ggml_backend_tensor_set(l.computed_wk_b.get(), wk_b->data, 0, ggml_nbytes(wk_b)); ggml_backend_tensor_set(l.computed_wk_b.get(), wk_b->data, 0, ggml_nbytes(wk_b));
if (ggml_backend_buffer_is_host(l.computed_wk_b->buffer)) {
iqk_modify_tensor(l.computed_wk_b.get());
}
l.wk_b = l.computed_wk_b.get(); l.wk_b = l.computed_wk_b.get();
@ -8243,6 +8247,9 @@ static bool llm_load_tensors(
ggml_set_name(l.computed_wv_b.get(), name.c_str()); ggml_set_name(l.computed_wv_b.get(), name.c_str());
ggml_backend_buffer_set_usage(l.computed_wv_b->buffer, GGML_BACKEND_BUFFER_USAGE_WEIGHTS); ggml_backend_buffer_set_usage(l.computed_wv_b->buffer, GGML_BACKEND_BUFFER_USAGE_WEIGHTS);
ggml_backend_tensor_set(l.computed_wv_b.get(), wv_b->data, 0, ggml_nbytes(wv_b)); ggml_backend_tensor_set(l.computed_wv_b.get(), wv_b->data, 0, ggml_nbytes(wv_b));
if (ggml_backend_buffer_is_host(l.computed_wv_b->buffer)) {
iqk_modify_tensor(l.computed_wv_b.get());
}
l.wv_b = l.computed_wv_b.get(); l.wv_b = l.computed_wv_b.get();
@ -17140,11 +17147,48 @@ static size_t llama_tensor_quantize_internal(enum ggml_type new_type, const floa
return new_size; return new_size;
} }
static llama_ftype repacked_ftype(llama_ftype ftype) {
static std::unordered_map<llama_ftype, llama_ftype> k_map = {
{ LLAMA_FTYPE_MOSTLY_Q4_0, LLAMA_FTYPE_MOSTLY_Q4_0_R8 },
{ LLAMA_FTYPE_MOSTLY_Q8_0, LLAMA_FTYPE_MOSTLY_Q8_0_R8 },
{ LLAMA_FTYPE_MOSTLY_Q5_0, LLAMA_FTYPE_MOSTLY_Q5_0_R4 },
{ LLAMA_FTYPE_MOSTLY_Q2_K, LLAMA_FTYPE_MOSTLY_Q2_K_R4 },
{ LLAMA_FTYPE_MOSTLY_Q3_K_S, LLAMA_FTYPE_MOSTLY_Q3_K_R4 },
{ LLAMA_FTYPE_MOSTLY_Q3_K_M, LLAMA_FTYPE_MOSTLY_Q3_K_R4 },
{ LLAMA_FTYPE_MOSTLY_Q3_K_L, LLAMA_FTYPE_MOSTLY_Q3_K_R4 },
{ LLAMA_FTYPE_MOSTLY_Q4_K_S, LLAMA_FTYPE_MOSTLY_Q4_K_R4 },
{ LLAMA_FTYPE_MOSTLY_Q4_K_M, LLAMA_FTYPE_MOSTLY_Q4_K_R4 },
{ LLAMA_FTYPE_MOSTLY_Q5_K_S, LLAMA_FTYPE_MOSTLY_Q5_K_R4 },
{ LLAMA_FTYPE_MOSTLY_Q5_K_M, LLAMA_FTYPE_MOSTLY_Q5_K_R4 },
{ LLAMA_FTYPE_MOSTLY_Q6_K, LLAMA_FTYPE_MOSTLY_Q6_K_R4 },
{ LLAMA_FTYPE_MOSTLY_IQ2_XXS, LLAMA_FTYPE_MOSTLY_IQ2_XXS_R4 },
{ LLAMA_FTYPE_MOSTLY_IQ2_XS, LLAMA_FTYPE_MOSTLY_IQ2_XS_R4 },
{ LLAMA_FTYPE_MOSTLY_IQ3_XXS, LLAMA_FTYPE_MOSTLY_IQ3_XXS_R4 },
{ LLAMA_FTYPE_MOSTLY_IQ1_S, LLAMA_FTYPE_MOSTLY_IQ1_S_R4 },
{ LLAMA_FTYPE_MOSTLY_IQ4_NL, LLAMA_FTYPE_MOSTLY_IQ4_NL_R4 },
{ LLAMA_FTYPE_MOSTLY_IQ3_S, LLAMA_FTYPE_MOSTLY_IQ3_S_R4 },
{ LLAMA_FTYPE_MOSTLY_IQ2_M, LLAMA_FTYPE_MOSTLY_IQ2_M_R4 },
{ LLAMA_FTYPE_MOSTLY_IQ4_XS, LLAMA_FTYPE_MOSTLY_IQ4_XS_R8 },
{ LLAMA_FTYPE_MOSTLY_IQ1_M, LLAMA_FTYPE_MOSTLY_IQ1_M_R4 },
{ LLAMA_FTYPE_MOSTLY_Q6_0, LLAMA_FTYPE_MOSTLY_Q6_0_R4 },
{ LLAMA_FTYPE_MOSTLY_BF16, LLAMA_FTYPE_MOSTLY_BF16_R16 },
{ LLAMA_FTYPE_MOSTLY_IQ2_BN, LLAMA_FTYPE_MOSTLY_IQ2_BN_R4 },
{ LLAMA_FTYPE_MOSTLY_IQ2_K, LLAMA_FTYPE_MOSTLY_IQ2_K_R4 },
{ LLAMA_FTYPE_MOSTLY_IQ3_K, LLAMA_FTYPE_MOSTLY_IQ3_K_R4 },
{ LLAMA_FTYPE_MOSTLY_IQ4_K, LLAMA_FTYPE_MOSTLY_IQ4_K_R4 },
{ LLAMA_FTYPE_MOSTLY_IQ5_K, LLAMA_FTYPE_MOSTLY_IQ5_K_R4 },
{ LLAMA_FTYPE_MOSTLY_IQ4_KS, LLAMA_FTYPE_MOSTLY_IQ4_KS_R4 },
{ LLAMA_FTYPE_MOSTLY_Q8_KV, LLAMA_FTYPE_MOSTLY_Q8_KV_R8 },
};
if (auto it = k_map.find(ftype); it != k_map.end()) return it->second;
return ftype;
}
static void llama_model_quantize_internal(const std::string & fname_inp, const std::string & fname_out, const llama_model_quantize_params * params) { static void llama_model_quantize_internal(const std::string & fname_inp, const std::string & fname_out, const llama_model_quantize_params * params) {
ggml_type default_type; ggml_type default_type;
llama_ftype ftype = params->ftype; llama_ftype ftype = params->ftype;
switch (params->ftype) { switch (ftype) {
case LLAMA_FTYPE_MOSTLY_Q4_0: default_type = GGML_TYPE_Q4_0; break; case LLAMA_FTYPE_MOSTLY_Q4_0: default_type = GGML_TYPE_Q4_0; break;
case LLAMA_FTYPE_MOSTLY_Q4_1: default_type = GGML_TYPE_Q4_1; break; case LLAMA_FTYPE_MOSTLY_Q4_1: default_type = GGML_TYPE_Q4_1; break;
case LLAMA_FTYPE_MOSTLY_Q5_0: default_type = GGML_TYPE_Q5_0; break; case LLAMA_FTYPE_MOSTLY_Q5_0: default_type = GGML_TYPE_Q5_0; break;
@ -17256,7 +17300,7 @@ static void llama_model_quantize_internal(const std::string & fname_inp, const s
ftype = model.ftype; ftype = model.ftype;
} }
const std::unordered_map<std::string, std::vector<float>> * imatrix_data = nullptr; const std::unordered_map<std::string, std::vector<float>> * imatrix_data = nullptr;
if (params->imatrix) { if (!params->only_repack && params->imatrix) {
imatrix_data = static_cast<const std::unordered_map<std::string, std::vector<float>>*>(params->imatrix); imatrix_data = static_cast<const std::unordered_map<std::string, std::vector<float>>*>(params->imatrix);
if (imatrix_data) { if (imatrix_data) {
LLAMA_LOG_INFO("================================ Have weights data with %d entries\n",int(imatrix_data->size())); LLAMA_LOG_INFO("================================ Have weights data with %d entries\n",int(imatrix_data->size()));
@ -17278,7 +17322,6 @@ static void llama_model_quantize_internal(const std::string & fname_inp, const s
// copy the KV pairs from the input file // copy the KV pairs from the input file
gguf_set_kv (ctx_out, ml.meta); gguf_set_kv (ctx_out, ml.meta);
gguf_set_val_u32(ctx_out, "general.quantization_version", GGML_QNT_VERSION); // TODO: use LLM_KV gguf_set_val_u32(ctx_out, "general.quantization_version", GGML_QNT_VERSION); // TODO: use LLM_KV
gguf_set_val_u32(ctx_out, "general.file_type", ftype); // TODO: use LLM_KV
// Remove split metadata // Remove split metadata
gguf_remove_key(ctx_out, ml.llm_kv(LLM_KV_SPLIT_NO).c_str()); gguf_remove_key(ctx_out, ml.llm_kv(LLM_KV_SPLIT_NO).c_str());
@ -17303,9 +17346,20 @@ static void llama_model_quantize_internal(const std::string & fname_inp, const s
} }
} }
bool is_repacked = ml.ftype >= LLAMA_FTYPE_MOSTLY_Q4_0_R8 && ml.ftype <= LLAMA_FTYPE_MOSTLY_Q8_K_R8;
int n_to_repack = 0, n_to_modify = 0;
for (int i = 0; i < ml.n_tensors; ++i) { for (int i = 0; i < ml.n_tensors; ++i) {
const struct ggml_tensor * meta = ml.get_tensor_meta(i); const struct ggml_tensor * meta = ml.get_tensor_meta(i);
if (params->only_repack) {
auto repacked_type = (ggml_type)iqk_repacked_type(meta);
if (repacked_type != meta->type) {
++n_to_repack;
} else if (!is_repacked) {
if (iqk_should_modify_tensor(meta)) ++n_to_modify;
}
}
const std::string name = ggml_get_name(meta); const std::string name = ggml_get_name(meta);
// TODO: avoid hardcoded tensor names - use the TN_* constants // TODO: avoid hardcoded tensor names - use the TN_* constants
@ -17317,6 +17371,18 @@ static void llama_model_quantize_internal(const std::string & fname_inp, const s
} }
} }
if (params->only_repack) {
if (n_to_repack == 0 && n_to_modify == 0) {
printf("=========================== %s: nothing to do for only_repack option\n", __func__);
return;
}
ftype = repacked_ftype(model.ftype);
printf("===================== Model ftype: %s: Repacked ftype: %s\n", llama_model_ftype_name(model.ftype).c_str(),
llama_model_ftype_name(ftype).c_str());
}
gguf_set_val_u32(ctx_out, "general.file_type", ftype); // TODO: use LLM_KV
qs.n_ffn_down = qs.n_ffn_gate = qs.n_ffn_up = (int)model.hparams.n_layer; qs.n_ffn_down = qs.n_ffn_gate = qs.n_ffn_up = (int)model.hparams.n_layer;
// sanity checks // sanity checks
@ -17457,6 +17523,36 @@ static void llama_model_quantize_internal(const std::string & fname_inp, const s
void * new_data; void * new_data;
size_t new_size; size_t new_size;
if (params->only_repack) {
ggml_type repacked_type = (ggml_type)iqk_repacked_type(tensor);
bool modify = !is_repacked && iqk_should_modify_tensor(tensor);
if (modify || repacked_type != tensor->type) {
new_type = repacked_type;
new_size = ggml_nbytes(tensor);
if ((int)work.size() < new_size) work.resize(new_size);
new_data = work.data();
auto aux_tensor = *tensor;
aux_tensor.data = work.data();
std::memcpy(aux_tensor.data, tensor->data, new_size);
if (repacked_type != tensor->type) {
iqk_repack_tensor(&aux_tensor);
GGML_ASSERT(aux_tensor.type == repacked_type);
} else {
bool did_modify = iqk_modify_tensor(&aux_tensor);
GGML_ASSERT(did_modify);
}
}
else {
new_type = tensor->type;
new_size = ggml_nbytes(tensor);
new_data = tensor->data;
}
LLAMA_LOG_INFO("size = %8.3f MB, type = %s\n", new_size/1024.0/1024.0, ggml_type_name(new_type));
goto QuantizationDone;
}
if (quantize) { if (quantize) {
new_type = default_type; new_type = default_type;
if (new_type == GGML_TYPE_BF16_R16 && strcmp(tensor->name, "token_embd.weight") == 0) { if (new_type == GGML_TYPE_BF16_R16 && strcmp(tensor->name, "token_embd.weight") == 0) {
@ -17562,7 +17658,7 @@ static void llama_model_quantize_internal(const std::string & fname_inp, const s
new_type == GGML_TYPE_IQ1_S_R4|| new_type == GGML_TYPE_IQ1_S_R4||
new_type == GGML_TYPE_IQ1_M_R4|| new_type == GGML_TYPE_IQ1_M_R4||
(new_type == GGML_TYPE_IQ1_M && strcmp(tensor->name, "token_embd.weight") && strcmp(tensor->name, "output.weight")) || (new_type == GGML_TYPE_IQ1_M && strcmp(tensor->name, "token_embd.weight") && strcmp(tensor->name, "output.weight")) ||
(new_type == GGML_TYPE_Q2_K && params->ftype == LLAMA_FTYPE_MOSTLY_Q2_K_S && strcmp(tensor->name, "token_embd.weight") != 0))) { (new_type == GGML_TYPE_Q2_K && ftype == LLAMA_FTYPE_MOSTLY_Q2_K_S && strcmp(tensor->name, "token_embd.weight") != 0))) {
LLAMA_LOG_ERROR("\n\n============================================================\n"); LLAMA_LOG_ERROR("\n\n============================================================\n");
LLAMA_LOG_ERROR("Missing importance matrix for tensor %s in a very low-bit quantization\n", tensor->name); LLAMA_LOG_ERROR("Missing importance matrix for tensor %s in a very low-bit quantization\n", tensor->name);
LLAMA_LOG_ERROR("The result will be garbage, so bailing out\n"); LLAMA_LOG_ERROR("The result will be garbage, so bailing out\n");
@ -17727,6 +17823,8 @@ static void llama_model_quantize_internal(const std::string & fname_inp, const s
} }
LLAMA_LOG_INFO("size = %8.2f MiB -> %8.2f MiB\n", ggml_nbytes(tensor)/1024.0/1024.0, new_size/1024.0/1024.0); LLAMA_LOG_INFO("size = %8.2f MiB -> %8.2f MiB\n", ggml_nbytes(tensor)/1024.0/1024.0, new_size/1024.0/1024.0);
} }
QuantizationDone:;
total_size_org += ggml_nbytes(tensor); total_size_org += ggml_nbytes(tensor);
total_size_new += new_size; total_size_new += new_size;
@ -18051,6 +18149,7 @@ struct llama_model_quantize_params llama_model_quantize_default_params() {
/*.pure =*/ false, /*.pure =*/ false,
/*.keep_split =*/ false, /*.keep_split =*/ false,
/*.ignore_imatrix_rules =*/ false, /*.ignore_imatrix_rules =*/ false,
/*.only_repack =*/ false,
/*.imatrix =*/ nullptr, /*.imatrix =*/ nullptr,
/*.kv_overrides =*/ nullptr, /*.kv_overrides =*/ nullptr,
/*.custom_quants =*/ nullptr, /*.custom_quants =*/ nullptr,