From f6602bf0e049332708d38fbd4e6522dc60c24079 Mon Sep 17 00:00:00 2001 From: Ed Addario Date: Tue, 30 Jun 2026 17:37:43 +0100 Subject: [PATCH 01/21] Add IQ2_NL and IQ3_NL types --- ggml/include/ggml.h | 4 +++- ggml/src/ggml-cpu/ops.cpp | 2 ++ gguf-py/gguf/constants.py | 2 ++ 3 files changed, 7 insertions(+), 1 deletion(-) diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h index d6807b6dd47..f62eff49913 100644 --- a/ggml/include/ggml.h +++ b/ggml/include/ggml.h @@ -429,7 +429,9 @@ extern "C" { GGML_TYPE_MXFP4 = 39, // MXFP4 (1 block) GGML_TYPE_NVFP4 = 40, // NVFP4 (4 blocks, E4M3 scale) GGML_TYPE_Q1_0 = 41, - GGML_TYPE_COUNT = 42, + GGML_TYPE_IQ2_NL = 42, + GGML_TYPE_IQ3_NL = 43, + GGML_TYPE_COUNT = 44, }; // precision diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index 6724686b8ae..faf92ca174f 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -5692,6 +5692,8 @@ void ggml_compute_forward_clamp( case GGML_TYPE_IQ4_XS: case GGML_TYPE_IQ3_S: case GGML_TYPE_IQ2_S: + case GGML_TYPE_IQ2_NL: + case GGML_TYPE_IQ3_NL: case GGML_TYPE_Q8_K: case GGML_TYPE_I8: case GGML_TYPE_I16: diff --git a/gguf-py/gguf/constants.py b/gguf-py/gguf/constants.py index 1bda9452dde..1e5a307234b 100644 --- a/gguf-py/gguf/constants.py +++ b/gguf-py/gguf/constants.py @@ -4415,6 +4415,8 @@ class GGMLQuantizationType(IntEnum): MXFP4 = 39 NVFP4 = 40 Q1_0 = 41 + IQ2_NL = 42 + IQ3_NL = 43 class ExpertGatingFuncType(IntEnum): From a6c99f80141aec4bf77a0e4a7a31ec948760d5f2 Mon Sep 17 00:00:00 2001 From: Ed Addario Date: Tue, 30 Jun 2026 17:45:39 +0100 Subject: [PATCH 02/21] Add block structs and codebook tables --- ggml/src/ggml-common.h | 23 +++++++++++++++++++++++ 1 file changed, 23 insertions(+) diff --git a/ggml/src/ggml-common.h b/ggml/src/ggml-common.h index f05683b44cd..68f58a55e51 100644 --- a/ggml/src/ggml-common.h +++ b/ggml/src/ggml-common.h @@ -441,6 +441,21 @@ typedef struct { } block_iq4_nl; static_assert(sizeof(block_iq4_nl) == sizeof(ggml_half) + QK4_NL/2, "wrong iq4_nl block size/padding"); +#define QK2_NL 32 +typedef struct { + ggml_half d; + uint8_t qs[QK2_NL/4]; +} block_iq2_nl; +static_assert(sizeof(block_iq2_nl) == sizeof(ggml_half) + QK2_NL/4, "wrong iq2_nl block size/padding"); + +#define QK3_NL 32 +typedef struct { + ggml_half d; + uint8_t qh[QK3_NL/8]; + uint8_t qs[QK3_NL/4]; +} block_iq3_nl; +static_assert(sizeof(block_iq3_nl) == sizeof(ggml_half) + QK3_NL/8 + QK3_NL/4, "wrong iq3_nl block size/padding"); + typedef struct { ggml_half d; uint16_t scales_h; @@ -1111,6 +1126,14 @@ GGML_TABLE_BEGIN(int8_t, kvalues_iq4nl, 16) -127, -104, -83, -65, -49, -35, -22, -10, 1, 13, 25, 38, 53, 69, 89, 113, GGML_TABLE_END() +GGML_TABLE_BEGIN(int8_t, kvalues_iq2nl, 4) + -127, -38, 38, 127, +GGML_TABLE_END() + +GGML_TABLE_BEGIN(int8_t, kvalues_iq3nl, 8) + -127, -79, -45, -14, 14, 45, 79, 127, +GGML_TABLE_END() + // e2m1 values (doubled) // ref: https://www.opencompute.org/documents/ocp-microscaling-formats-mx-v1-0-spec-final-pdf GGML_TABLE_BEGIN(int8_t, kvalues_mxfp4, 16) From 427d20ec333f7ff70857c81f2349c5c19757e931 Mon Sep 17 00:00:00 2001 From: Ed Addario Date: Tue, 30 Jun 2026 17:50:55 +0100 Subject: [PATCH 03/21] Mirror IQ4_NL --- ggml/src/ggml-quants.h | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/ggml/src/ggml-quants.h b/ggml/src/ggml-quants.h index d56c86da890..4889a6a53af 100644 --- a/ggml/src/ggml-quants.h +++ b/ggml/src/ggml-quants.h @@ -38,6 +38,8 @@ GGML_API void quantize_row_tq2_0_ref(const float * GGML_RESTRICT x, block_tq2_0 GGML_API void quantize_row_iq3_xxs_ref(const float * GGML_RESTRICT x, block_iq3_xxs * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_iq4_nl_ref (const float * GGML_RESTRICT x, block_iq4_nl * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_iq4_xs_ref (const float * GGML_RESTRICT x, block_iq4_xs * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_iq2_nl_ref (const float * GGML_RESTRICT x, block_iq2_nl * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_iq3_nl_ref (const float * GGML_RESTRICT x, block_iq3_nl * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_iq3_s_ref (const float * GGML_RESTRICT x, block_iq3_s * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_iq2_s_ref (const float * GGML_RESTRICT x, block_iq2_s * GGML_RESTRICT y, int64_t k); @@ -71,6 +73,8 @@ GGML_API void dequantize_row_iq1_s (const block_iq1_s * GGML_RESTRICT x, floa GGML_API void dequantize_row_iq1_m (const block_iq1_m * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_iq4_nl (const block_iq4_nl * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_iq4_xs (const block_iq4_xs * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_iq2_nl (const block_iq2_nl * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_iq3_nl (const block_iq3_nl * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_iq3_s (const block_iq3_s * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); // Quantization utilizing an importance matrix (a.k.a. "Activation aWare Quantization") @@ -82,6 +86,8 @@ GGML_API size_t quantize_iq1_s (const float * GGML_RESTRICT src, void * GGML_RE GGML_API size_t quantize_iq1_m (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_iq4_nl (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_iq4_xs (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); +GGML_API size_t quantize_iq2_nl (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); +GGML_API size_t quantize_iq3_nl (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_iq3_s (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_tq1_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); From da119d8aa8985787958c761581e1e61db9a84e7b Mon Sep 17 00:00:00 2001 From: Ed Addario Date: Tue, 30 Jun 2026 22:22:40 +0100 Subject: [PATCH 04/21] Add quantize / dequantize logic --- ggml/src/ggml-quants.c | 167 +++++++++++++++++++++++++++++++++++++---- 1 file changed, 151 insertions(+), 16 deletions(-) diff --git a/ggml/src/ggml-quants.c b/ggml/src/ggml-quants.c index 15d231f70c0..c3d649dafc7 100644 --- a/ggml/src/ggml-quants.c +++ b/ggml/src/ggml-quants.c @@ -2668,6 +2668,43 @@ void dequantize_row_iq4_nl(const block_iq4_nl * GGML_RESTRICT x, float * GGML_RE } } +void dequantize_row_iq2_nl(const block_iq2_nl * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + assert(k % QK2_NL == 0); + const int64_t nb = k / QK2_NL; + + for (int i = 0; i < nb; i++) { + const uint8_t * qs = x[i].qs; + + const float d = GGML_FP16_TO_FP32(x[i].d); + for (int j = 0; j < QK2_NL/4; ++j) { + y[j + 0*(QK2_NL/4)] = d * kvalues_iq2nl[(qs[j] >> 0) & 3]; + y[j + 1*(QK2_NL/4)] = d * kvalues_iq2nl[(qs[j] >> 2) & 3]; + y[j + 2*(QK2_NL/4)] = d * kvalues_iq2nl[(qs[j] >> 4) & 3]; + y[j + 3*(QK2_NL/4)] = d * kvalues_iq2nl[(qs[j] >> 6) & 3]; + } + y += QK2_NL; + } +} + +void dequantize_row_iq3_nl(const block_iq3_nl * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + assert(k % QK3_NL == 0); + const int64_t nb = k / QK3_NL; + + for (int i = 0; i < nb; i++) { + const uint8_t * qs = x[i].qs; + const uint8_t * qh = x[i].qh; + + const float d = GGML_FP16_TO_FP32(x[i].d); + for (int j = 0; j < QK3_NL/4; ++j) { + for (int g = 0; g < 4; ++g) { + const int idx = ((qs[j] >> (2*g)) & 3) | (((qh[g] >> j) & 1) << 2); + y[j + g*(QK3_NL/4)] = d * kvalues_iq3nl[idx]; + } + } + y += QK3_NL; + } +} + void dequantize_row_iq4_xs(const block_iq4_xs * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { assert(k % QK_K == 0); const int64_t nb = k / QK_K; @@ -4891,8 +4928,9 @@ size_t quantize_iq1_m(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, // ============================ 4-bit non-linear quants -static void quantize_row_iq4_nl_impl(const int super_block_size, const int block_size, const float * GGML_RESTRICT x, - ggml_fp16_t * dh, uint8_t * q4, uint16_t * scales_h, uint8_t * scales_l, +static void quantize_row_iq4_nl_impl(const int super_block_size, const int block_size, const int nvalues, + const float * GGML_RESTRICT x, + ggml_fp16_t * dh, uint8_t * q4, uint8_t * qh, uint16_t * scales_h, uint8_t * scales_l, float * scales, float * weight, uint8_t * L, const int8_t * values, const float * quant_weights, @@ -4902,7 +4940,8 @@ static void quantize_row_iq4_nl_impl(const int super_block_size, const int block for (int j = 0; j < super_block_size; ++j) sigma2 += x[j]*x[j]; sigma2 *= 2.f/super_block_size; - memset(q4, 0, super_block_size/2); + memset(q4, 0, nvalues > 8 ? super_block_size/2 : super_block_size/4); + if (qh) memset(qh, 0, super_block_size/8); dh[0] = GGML_FP32_TO_FP16(0.f); float max_scale = 0, amax_scale = 0; @@ -4931,7 +4970,7 @@ static void quantize_row_iq4_nl_impl(const int super_block_size, const int block float sumqx = 0, sumq2 = 0; for (int j = 0; j < block_size; ++j) { float al = id*xb[j]; - int l = best_index_int8(16, values, al); + int l = best_index_int8(nvalues, values, al); Lb[j] = l; float q = values[l]; float w = weight[j]; @@ -4945,7 +4984,7 @@ static void quantize_row_iq4_nl_impl(const int super_block_size, const int block sumqx = sumq2 = 0; for (int j = 0; j < block_size; ++j) { float al = id*xb[j]; - int l = best_index_int8(16, values, al); + int l = best_index_int8(nvalues, values, al); float q = values[l]; float w = weight[j]; sumqx += w*q*xb[j]; @@ -4976,7 +5015,7 @@ static void quantize_row_iq4_nl_impl(const int super_block_size, const int block uint8_t * Lb = L + ib*block_size; const float * xb = x + ib*block_size; for (int j = 0; j < block_size; ++j) { - Lb[j] = best_index_int8(16, values, idl*xb[j]); + Lb[j] = best_index_int8(nvalues, values, idl*xb[j]); } l += 32; uint8_t l_l = l & 0xf; @@ -4990,14 +5029,28 @@ static void quantize_row_iq4_nl_impl(const int super_block_size, const int block if (ntry > 0) { float id = scales[0] ? 1/scales[0] : 0; for (int j = 0; j < super_block_size; ++j) { - L[j] = best_index_int8(16, values, id*x[j]); + L[j] = best_index_int8(nvalues, values, id*x[j]); } } } - for (int i = 0; i < super_block_size/32; ++i) { - for (int j = 0; j < 16; ++j) { - q4[16*i + j] = L[32*i + j] | (L[32*i + 16 + j] << 4); + if (nvalues > 8) { + for (int i = 0; i < super_block_size/32; ++i) { + for (int j = 0; j < 16; ++j) { + q4[16*i + j] = L[32*i + j] | (L[32*i + 16 + j] << 4); + } + } + } else { + const int n = super_block_size/4; + for (int j = 0; j < n; ++j) { + q4[j] = (L[j] & 3) | ((L[j+n] & 3) << 2) | ((L[j+2*n] & 3) << 4) | ((L[j+3*n] & 3) << 6); + } + if (qh) { + for (int g = 0; g < 4; ++g) { + uint8_t h = 0; + for (int j = 0; j < n; ++j) h |= ((L[j+g*n] >> 2) & 1) << j; + qh[g] = h; + } } } } @@ -5015,8 +5068,8 @@ size_t quantize_iq4_nl(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst block_iq4_nl * iq4 = (block_iq4_nl *)qrow; for (int ibl = 0; ibl < nblock; ++ibl) { const float * qw = quant_weights ? quant_weights + QK4_NL*ibl : NULL; - quantize_row_iq4_nl_impl(QK4_NL, 32, src + QK4_NL*ibl, &iq4[ibl].d, iq4[ibl].qs, &unused_h, unused_l, - &scale, weight, L, kvalues_iq4nl, qw, 7); + quantize_row_iq4_nl_impl(QK4_NL, 32, 16, src + QK4_NL*ibl, &iq4[ibl].d, iq4[ibl].qs, NULL, &unused_h, + unused_l, &scale, weight, L, kvalues_iq4nl, qw, 7); } src += n_per_row; qrow += nblock*sizeof(block_iq4_nl); @@ -5035,8 +5088,8 @@ void quantize_row_iq4_nl_ref(const float * GGML_RESTRICT x, block_iq4_nl * GGML_ float scale; block_iq4_nl * iq4 = y; for (int ibl = 0; ibl < nblock; ++ibl) { - quantize_row_iq4_nl_impl(QK4_NL, 32, x + QK4_NL*ibl, &iq4[ibl].d, iq4[ibl].qs, &unused_h, unused_l, - &scale, weight, L, kvalues_iq4nl, NULL, -1); + quantize_row_iq4_nl_impl(QK4_NL, 32, 16, x + QK4_NL*ibl, &iq4[ibl].d, iq4[ibl].qs, NULL, &unused_h, + unused_l, &scale, weight, L, kvalues_iq4nl, NULL, -1); } } @@ -5051,8 +5104,8 @@ size_t quantize_iq4_xs(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst block_iq4_xs * iq4 = (block_iq4_xs *)qrow; for (int ibl = 0; ibl < nblock; ++ibl) { const float * qw = quant_weights ? quant_weights + QK_K*ibl : NULL; - quantize_row_iq4_nl_impl(QK_K, 32, src + QK_K*ibl, &iq4[ibl].d, iq4[ibl].qs, &iq4[ibl].scales_h, iq4[ibl].scales_l, - scales, weight, L, kvalues_iq4nl, qw, 7); + quantize_row_iq4_nl_impl(QK_K, 32, 16, src + QK_K*ibl, &iq4[ibl].d, iq4[ibl].qs, NULL, + &iq4[ibl].scales_h, iq4[ibl].scales_l, scales, weight, L, kvalues_iq4nl, qw, 7); } src += n_per_row; qrow += nblock*sizeof(block_iq4_xs); @@ -5065,6 +5118,80 @@ void quantize_row_iq4_xs_ref(const float * GGML_RESTRICT x, block_iq4_xs * GGML_ quantize_iq4_xs(x, y, 1, k, NULL); } +size_t quantize_iq2_nl(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { + GGML_ASSERT(n_per_row%QK2_NL == 0); + int64_t nblock = n_per_row/QK2_NL; + char * qrow = (char *)dst; + uint8_t L[QK2_NL]; + float weight[QK2_NL]; + uint16_t unused_h; + uint8_t * unused_l = NULL; + float scale; + for (int64_t row = 0; row < nrow; ++row) { + block_iq2_nl * iq2 = (block_iq2_nl *)qrow; + for (int ibl = 0; ibl < nblock; ++ibl) { + const float * qw = quant_weights ? quant_weights + QK2_NL*ibl : NULL; + quantize_row_iq4_nl_impl(QK2_NL, 32, 4, src + QK2_NL*ibl, &iq2[ibl].d, iq2[ibl].qs, NULL, + &unused_h, unused_l, &scale, weight, L, kvalues_iq2nl, qw, 7); + } + src += n_per_row; + qrow += nblock*sizeof(block_iq2_nl); + } + return nrow * nblock * sizeof(block_iq2_nl); +} + +void quantize_row_iq2_nl_ref(const float * GGML_RESTRICT x, block_iq2_nl * GGML_RESTRICT y, int64_t k) { + GGML_ASSERT(k%QK2_NL == 0); + int64_t nblock = k/QK2_NL; + uint8_t L[QK2_NL]; + float weight[QK2_NL]; + uint16_t unused_h; + uint8_t * unused_l = NULL; + float scale; + block_iq2_nl * iq2 = y; + for (int ibl = 0; ibl < nblock; ++ibl) { + quantize_row_iq4_nl_impl(QK2_NL, 32, 4, x + QK2_NL*ibl, &iq2[ibl].d, iq2[ibl].qs, NULL, + &unused_h, unused_l, &scale, weight, L, kvalues_iq2nl, NULL, -1); + } +} + +size_t quantize_iq3_nl(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { + GGML_ASSERT(n_per_row%QK3_NL == 0); + int64_t nblock = n_per_row/QK3_NL; + char * qrow = (char *)dst; + uint8_t L[QK3_NL]; + float weight[QK3_NL]; + uint16_t unused_h; + uint8_t * unused_l = NULL; + float scale; + for (int64_t row = 0; row < nrow; ++row) { + block_iq3_nl * iq3 = (block_iq3_nl *)qrow; + for (int ibl = 0; ibl < nblock; ++ibl) { + const float * qw = quant_weights ? quant_weights + QK3_NL*ibl : NULL; + quantize_row_iq4_nl_impl(QK3_NL, 32, 8, src + QK3_NL*ibl, &iq3[ibl].d, iq3[ibl].qs, iq3[ibl].qh, + &unused_h, unused_l, &scale, weight, L, kvalues_iq3nl, qw, 7); + } + src += n_per_row; + qrow += nblock*sizeof(block_iq3_nl); + } + return nrow * nblock * sizeof(block_iq3_nl); +} + +void quantize_row_iq3_nl_ref(const float * GGML_RESTRICT x, block_iq3_nl * GGML_RESTRICT y, int64_t k) { + GGML_ASSERT(k%QK3_NL == 0); + int64_t nblock = k/QK3_NL; + uint8_t L[QK3_NL]; + float weight[QK3_NL]; + uint16_t unused_h; + uint8_t * unused_l = NULL; + float scale; + block_iq3_nl * iq3 = y; + for (int ibl = 0; ibl < nblock; ++ibl) { + quantize_row_iq4_nl_impl(QK3_NL, 32, 8, x + QK3_NL*ibl, &iq3[ibl].d, iq3[ibl].qs, iq3[ibl].qh, + &unused_h, unused_l, &scale, weight, L, kvalues_iq3nl, NULL, -1); + } +} + // =============================== 2.5625 bpw static void quantize_row_iq2_s_impl(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t n, const float * GGML_RESTRICT quant_weights) { @@ -5573,6 +5700,14 @@ bool ggml_validate_row_data(enum ggml_type type, const void * data, size_t nbyte { VALIDATE_ROW_DATA_D_F16_IMPL(block_iq4_nl, data, nb); } break; + case GGML_TYPE_IQ2_NL: + { + VALIDATE_ROW_DATA_D_F16_IMPL(block_iq2_nl, data, nb); + } break; + case GGML_TYPE_IQ3_NL: + { + VALIDATE_ROW_DATA_D_F16_IMPL(block_iq3_nl, data, nb); + } break; case GGML_TYPE_I8: case GGML_TYPE_I16: From f4afcfc8601e4e19accfc7dc8bb3a80f1fc901d3 Mon Sep 17 00:00:00 2001 From: Ed Addario Date: Tue, 30 Jun 2026 22:27:16 +0100 Subject: [PATCH 05/21] Add type traits --- ggml/src/ggml.c | 18 ++++++++++++++++++ 1 file changed, 18 insertions(+) diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index 0f682fd1856..bf3f0491726 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -842,6 +842,22 @@ static const struct ggml_type_traits type_traits[GGML_TYPE_COUNT] = { .to_float = (ggml_to_float_t) dequantize_row_iq1_m, .from_float_ref = NULL, }, + [GGML_TYPE_IQ2_NL] = { + .type_name = "iq2_nl", + .blck_size = QK2_NL, + .type_size = sizeof(block_iq2_nl), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_iq2_nl, + .from_float_ref = (ggml_from_float_t)quantize_row_iq2_nl_ref, + }, + [GGML_TYPE_IQ3_NL] = { + .type_name = "iq3_nl", + .blck_size = QK3_NL, + .type_size = sizeof(block_iq3_nl), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_iq3_nl, + .from_float_ref = (ggml_from_float_t)quantize_row_iq3_nl_ref, + }, [GGML_TYPE_IQ4_NL] = { .type_name = "iq4_nl", .blck_size = QK4_NL, @@ -7750,6 +7766,8 @@ size_t ggml_quantize_chunk( case GGML_TYPE_IQ2_S: result = quantize_iq2_s (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_IQ1_S: result = quantize_iq1_s (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_IQ1_M: result = quantize_iq1_m (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_IQ2_NL: result = quantize_iq2_nl (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_IQ3_NL: result = quantize_iq3_nl (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_XS: result = quantize_iq4_xs (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_F16: From c3ad8a480f3f857c2a59bd4da8a361e3b3badcb7 Mon Sep 17 00:00:00 2001 From: Ed Addario Date: Tue, 30 Jun 2026 22:30:06 +0100 Subject: [PATCH 06/21] Add definitions --- gguf-py/gguf/constants.py | 2 ++ gguf-py/gguf/quants.py | 46 +++++++++++++++++++++++++++++++++++++++ 2 files changed, 48 insertions(+) diff --git a/gguf-py/gguf/constants.py b/gguf-py/gguf/constants.py index 1e5a307234b..10e6fef5841 100644 --- a/gguf-py/gguf/constants.py +++ b/gguf-py/gguf/constants.py @@ -4596,6 +4596,8 @@ class VisionProjectorType: GGMLQuantizationType.MXFP4: (32, 1 + 16), GGMLQuantizationType.NVFP4: (64, 4 + 32), GGMLQuantizationType.Q1_0: (128, 2 + 16), + GGMLQuantizationType.IQ2_NL: (32, 2 + 8), + GGMLQuantizationType.IQ3_NL: (32, 2 + 4 + 8), } diff --git a/gguf-py/gguf/quants.py b/gguf-py/gguf/quants.py index 80966b6ef15..dd4ca627836 100644 --- a/gguf-py/gguf/quants.py +++ b/gguf-py/gguf/quants.py @@ -1328,6 +1328,52 @@ def dequantize_blocks(cls, blocks: np.ndarray) -> np.ndarray: return (dl * (grid + delta)).reshape((n_blocks, -1)) +class IQ2_NL(__Quant, qtype=GGMLQuantizationType.IQ2_NL): + kvalues = (-127, -38, 38, 127) + + @classmethod + def dequantize_blocks(cls, blocks: np.ndarray) -> np.ndarray: + n_blocks = blocks.shape[0] + + d, qs = np.hsplit(blocks, [2]) + + d = d.view(np.float16).astype(np.float32) + + qs = qs.reshape((n_blocks, -1, 1, cls.block_size // 4)) >> np.array([0, 2, 4, 6], dtype=np.uint8).reshape((1, 1, 4, 1)) + qs = (qs & np.uint8(0x03)).reshape((n_blocks, -1, 1)) + + kvalues = np.array(cls.kvalues, dtype=np.int8).reshape(1, 1, 4) + qs = np.take_along_axis(kvalues, qs, axis=-1).astype(np.float32).reshape((n_blocks, -1)) + + return (d * qs) + + +class IQ3_NL(__Quant, qtype=GGMLQuantizationType.IQ3_NL): + kvalues = (-127, -79, -45, -14, 14, 45, 79, 127) + + @classmethod + def dequantize_blocks(cls, blocks: np.ndarray) -> np.ndarray: + n_blocks = blocks.shape[0] + + d, rest = np.hsplit(blocks, [2]) + qh, qs = np.hsplit(rest, [cls.block_size // 8]) + + d = d.view(np.float16).astype(np.float32) + + lo = qs.reshape((n_blocks, -1, 1, cls.block_size // 4)) >> np.array([0, 2, 4, 6], dtype=np.uint8).reshape((1, 1, 4, 1)) + lo = (lo & np.uint8(0x03)).reshape((n_blocks, -1)) + + hi = qh.reshape((n_blocks, -1, 1)) >> np.array([0, 1, 2, 3, 4, 5, 6, 7], dtype=np.uint8).reshape((1, 1, 8)) + hi = (hi & np.uint8(0x01)).reshape((n_blocks, -1)) + + qs = (lo | (hi << np.uint8(2))).reshape((n_blocks, -1, 1)) + + kvalues = np.array(cls.kvalues, dtype=np.int8).reshape(1, 1, 8) + qs = np.take_along_axis(kvalues, qs, axis=-1).astype(np.float32).reshape((n_blocks, -1)) + + return (d * qs) + + class IQ4_NL(__Quant, qtype=GGMLQuantizationType.IQ4_NL): kvalues = (-127, -104, -83, -65, -49, -35, -22, -10, 1, 13, 25, 38, 53, 69, 89, 113) From 29191a8493f00d362707b4c87f143997f3cc2046 Mon Sep 17 00:00:00 2001 From: Ed Addario Date: Tue, 30 Jun 2026 22:45:29 +0100 Subject: [PATCH 07/21] Add generic ggml_vec_dot for iq2_nl and iq3_nl --- ggml/src/ggml-cpu/arch-fallback.h | 17 +++++++ ggml/src/ggml-cpu/quants.c | 73 +++++++++++++++++++++++++++++++ ggml/src/ggml-cpu/quants.h | 6 +++ 3 files changed, 96 insertions(+) diff --git a/ggml/src/ggml-cpu/arch-fallback.h b/ggml/src/ggml-cpu/arch-fallback.h index 1fc2b4b71bd..7e3ead33e2d 100644 --- a/ggml/src/ggml-cpu/arch-fallback.h +++ b/ggml/src/ggml-cpu/arch-fallback.h @@ -31,6 +31,8 @@ #define ggml_vec_dot_iq3_s_q8_K_generic ggml_vec_dot_iq3_s_q8_K #define ggml_vec_dot_iq1_s_q8_K_generic ggml_vec_dot_iq1_s_q8_K #define ggml_vec_dot_iq1_m_q8_K_generic ggml_vec_dot_iq1_m_q8_K +#define ggml_vec_dot_iq2_nl_q8_0_generic ggml_vec_dot_iq2_nl_q8_0 +#define ggml_vec_dot_iq3_nl_q8_0_generic ggml_vec_dot_iq3_nl_q8_0 #define ggml_vec_dot_iq4_nl_q8_0_generic ggml_vec_dot_iq4_nl_q8_0 #define ggml_vec_dot_iq4_xs_q8_K_generic ggml_vec_dot_iq4_xs_q8_K // repack.cpp @@ -71,6 +73,9 @@ #define ggml_gemm_q8_0_4x4_q8_0_generic ggml_gemm_q8_0_4x4_q8_0 #define ggml_gemm_q8_0_4x8_q8_0_generic ggml_gemm_q8_0_4x8_q8_0 #elif defined(__aarch64__) || defined(__arm__) || defined(_M_ARM) || defined(_M_ARM64) +// quants.c +#define ggml_vec_dot_iq2_nl_q8_0_generic ggml_vec_dot_iq2_nl_q8_0 +#define ggml_vec_dot_iq3_nl_q8_0_generic ggml_vec_dot_iq3_nl_q8_0 // repack.cpp #define ggml_quantize_mat_q8_K_4x4_generic ggml_quantize_mat_q8_K_4x4 #define ggml_quantize_mat_q8_K_4x8_generic ggml_quantize_mat_q8_K_4x8 @@ -83,6 +88,8 @@ #elif defined(__x86_64__) || defined(__i386__) || defined(_M_IX86) || defined(_M_X64) // quants.c #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0 +#define ggml_vec_dot_iq2_nl_q8_0_generic ggml_vec_dot_iq2_nl_q8_0 +#define ggml_vec_dot_iq3_nl_q8_0_generic ggml_vec_dot_iq3_nl_q8_0 // repack.cpp #define ggml_quantize_mat_q8_0_4x4_generic ggml_quantize_mat_q8_0_4x4 #define ggml_quantize_mat_q8_K_4x4_generic ggml_quantize_mat_q8_K_4x4 @@ -117,6 +124,8 @@ #define ggml_vec_dot_tq1_0_q8_K_generic ggml_vec_dot_tq1_0_q8_K #define ggml_vec_dot_tq2_0_q8_K_generic ggml_vec_dot_tq2_0_q8_K #define ggml_vec_dot_iq1_m_q8_K_generic ggml_vec_dot_iq1_m_q8_K +#define ggml_vec_dot_iq2_nl_q8_0_generic ggml_vec_dot_iq2_nl_q8_0 +#define ggml_vec_dot_iq3_nl_q8_0_generic ggml_vec_dot_iq3_nl_q8_0 // repack.cpp #define ggml_quantize_mat_q8_0_4x4_generic ggml_quantize_mat_q8_0_4x4 #define ggml_quantize_mat_q8_0_4x8_generic ggml_quantize_mat_q8_0_4x8 @@ -163,6 +172,8 @@ #define ggml_vec_dot_mxfp4_q8_0_generic ggml_vec_dot_mxfp4_q8_0 #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0 #define ggml_vec_dot_q1_0_q8_0_generic ggml_vec_dot_q1_0_q8_0 +#define ggml_vec_dot_iq2_nl_q8_0_generic ggml_vec_dot_iq2_nl_q8_0 +#define ggml_vec_dot_iq3_nl_q8_0_generic ggml_vec_dot_iq3_nl_q8_0 // repack.cpp #define ggml_quantize_mat_q8_0_4x4_generic ggml_quantize_mat_q8_0_4x4 #define ggml_quantize_mat_q8_0_4x8_generic ggml_quantize_mat_q8_0_4x8 @@ -203,6 +214,8 @@ #elif defined(__riscv) // quants.c #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0 +#define ggml_vec_dot_iq2_nl_q8_0_generic ggml_vec_dot_iq2_nl_q8_0 +#define ggml_vec_dot_iq3_nl_q8_0_generic ggml_vec_dot_iq3_nl_q8_0 // repack.cpp #define ggml_quantize_mat_q8_0_4x1_generic ggml_quantize_mat_q8_0_4x1 #define ggml_quantize_mat_q8_0_4x4_generic ggml_quantize_mat_q8_0_4x4 @@ -254,6 +267,8 @@ #define ggml_vec_dot_iq3_s_q8_K_generic ggml_vec_dot_iq3_s_q8_K #define ggml_vec_dot_iq1_s_q8_K_generic ggml_vec_dot_iq1_s_q8_K #define ggml_vec_dot_iq1_m_q8_K_generic ggml_vec_dot_iq1_m_q8_K +#define ggml_vec_dot_iq2_nl_q8_0_generic ggml_vec_dot_iq2_nl_q8_0 +#define ggml_vec_dot_iq3_nl_q8_0_generic ggml_vec_dot_iq3_nl_q8_0 // repack.cpp #define ggml_quantize_mat_q8_0_4x4_generic ggml_quantize_mat_q8_0_4x4 #define ggml_quantize_mat_q8_0_4x8_generic ggml_quantize_mat_q8_0_4x8 @@ -307,6 +322,8 @@ #define ggml_vec_dot_mxfp4_q8_0_generic ggml_vec_dot_mxfp4_q8_0 #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0 #define ggml_vec_dot_q1_0_q8_0_generic ggml_vec_dot_q1_0_q8_0 +#define ggml_vec_dot_iq2_nl_q8_0_generic ggml_vec_dot_iq2_nl_q8_0 +#define ggml_vec_dot_iq3_nl_q8_0_generic ggml_vec_dot_iq3_nl_q8_0 // repack.cpp #define ggml_quantize_mat_q8_0_4x4_generic ggml_quantize_mat_q8_0_4x4 #define ggml_quantize_mat_q8_0_4x8_generic ggml_quantize_mat_q8_0_4x8 diff --git a/ggml/src/ggml-cpu/quants.c b/ggml/src/ggml-cpu/quants.c index e5f9a4083f9..4473adf390a 100644 --- a/ggml/src/ggml-cpu/quants.c +++ b/ggml/src/ggml-cpu/quants.c @@ -1229,6 +1229,69 @@ void ggml_vec_dot_iq4_nl_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, *s = sumf; } +void ggml_vec_dot_iq2_nl_q8_0_generic(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) { + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + assert(n % QK2_NL == 0); + static_assert(QK2_NL == QK8_0, "QK2_NL and QK8_0 must be the same"); + + const block_iq2_nl * GGML_RESTRICT x = vx; + const block_q8_0 * GGML_RESTRICT y = vy; + + const int nb = n / QK2_NL; + + int ib = 0; + float sumf = 0; + + for (; ib < nb; ++ib) { + const float d = GGML_CPU_FP16_TO_FP32(y[ib].d)*GGML_CPU_FP16_TO_FP32(x[ib].d); + int sumi = 0; + for (int j = 0; j < QK2_NL/4; ++j) { + for (int g = 0; g < 4; ++g) { + sumi += y[ib].qs[j + g*(QK2_NL/4)] * kvalues_iq2nl[(x[ib].qs[j] >> 2*g) & 3]; + } + } + sumf += d * sumi; + } + *s = sumf; +} + +void ggml_vec_dot_iq3_nl_q8_0_generic(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) { + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + assert(n % QK3_NL == 0); + static_assert(QK3_NL == QK8_0, "QK3_NL and QK8_0 must be the same"); + + const block_iq3_nl * GGML_RESTRICT x = vx; + const block_q8_0 * GGML_RESTRICT y = vy; + + const int nb = n / QK3_NL; + + int ib = 0; + float sumf = 0; + + for (; ib < nb; ++ib) { + const float d = GGML_CPU_FP16_TO_FP32(y[ib].d)*GGML_CPU_FP16_TO_FP32(x[ib].d); + const uint8_t * qs = x[ib].qs; + const uint8_t * qh = x[ib].qh; + int sumi = 0; + for (int j = 0; j < QK3_NL/4; ++j) { + for (int g = 0; g < 4; ++g) { + const int idx = ((qs[j] >> 2*g) & 3) | (((qh[g] >> j) & 1) << 2); + sumi += y[ib].qs[j + g*(QK3_NL/4)] * kvalues_iq3nl[idx]; + } + } + sumf += d * sumi; + } + *s = sumf; +} + void ggml_vec_dot_iq4_xs_q8_K_generic(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) { assert(nrc == 1); UNUSED(nrc); @@ -1277,6 +1340,16 @@ void ggml_vec_dot_iq4_xs_q8_K_generic(int n, float * GGML_RESTRICT s, size_t bs, // ============================ 4-bit non-linear quants +void quantize_row_iq2_nl(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { + assert(k % QK2_NL == 0); + quantize_row_iq2_nl_ref(x, y, k); +} + +void quantize_row_iq3_nl(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { + assert(k % QK3_NL == 0); + quantize_row_iq3_nl_ref(x, y, k); +} + void quantize_row_iq4_nl(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { assert(k % QK4_NL == 0); quantize_row_iq4_nl_ref(x, y, k); diff --git a/ggml/src/ggml-cpu/quants.h b/ggml/src/ggml-cpu/quants.h index d4bc87a1c05..543612d3117 100644 --- a/ggml/src/ggml-cpu/quants.h +++ b/ggml/src/ggml-cpu/quants.h @@ -33,6 +33,8 @@ void quantize_row_q8_K(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, in void quantize_row_tq1_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_tq2_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_iq2_nl (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_iq3_nl (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_iq4_nl (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_iq4_xs (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); @@ -62,6 +64,8 @@ void ggml_vec_dot_iq2_s_q8_K (int n, float * GGML_RESTRICT s, size_t bs, const void ggml_vec_dot_iq3_xxs_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 ggml_vec_dot_iq1_s_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 ggml_vec_dot_iq1_m_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 ggml_vec_dot_iq2_nl_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 ggml_vec_dot_iq3_nl_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 ggml_vec_dot_iq4_nl_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 ggml_vec_dot_iq4_xs_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 ggml_vec_dot_iq3_s_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); @@ -95,6 +99,8 @@ void ggml_vec_dot_iq3_xxs_q8_K_generic(int n, float * GGML_RESTRICT s, size_t bs void ggml_vec_dot_iq3_s_q8_K_generic(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 ggml_vec_dot_iq1_s_q8_K_generic(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 ggml_vec_dot_iq1_m_q8_K_generic(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 ggml_vec_dot_iq2_nl_q8_0_generic(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 ggml_vec_dot_iq3_nl_q8_0_generic(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 ggml_vec_dot_iq4_nl_q8_0_generic(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 ggml_vec_dot_iq4_xs_q8_K_generic(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); From 6ea87768ec6ac906a06adad4b01179b9ec1cbb0c Mon Sep 17 00:00:00 2001 From: Ed Addario Date: Tue, 30 Jun 2026 22:48:54 +0100 Subject: [PATCH 08/21] Add type traits --- ggml/src/ggml-cpu/ggml-cpu.c | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c index eb8341c9aec..6d4d47df7c1 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -367,6 +367,18 @@ static const struct ggml_type_traits_cpu type_traits_cpu[GGML_TYPE_COUNT] = { .vec_dot_type = GGML_TYPE_Q8_K, .nrows = 1, }, + [GGML_TYPE_IQ2_NL] = { + .from_float = quantize_row_iq2_nl, + .vec_dot = ggml_vec_dot_iq2_nl_q8_0, + .vec_dot_type = GGML_TYPE_Q8_0, + .nrows = 1, + }, + [GGML_TYPE_IQ3_NL] = { + .from_float = quantize_row_iq3_nl, + .vec_dot = ggml_vec_dot_iq3_nl_q8_0, + .vec_dot_type = GGML_TYPE_Q8_0, + .nrows = 1, + }, [GGML_TYPE_IQ4_NL] = { .from_float = quantize_row_iq4_nl, .vec_dot = ggml_vec_dot_iq4_nl_q8_0, From 3d9f4ea688a8d54f361a47e91b827e31951a1bf3 Mon Sep 17 00:00:00 2001 From: Ed Addario Date: Tue, 30 Jun 2026 22:56:21 +0100 Subject: [PATCH 09/21] Add iq2_nl / iq3_nl --- ggml/src/ggml-cpu/ops.cpp | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index faf92ca174f..916efa41f82 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -684,6 +684,8 @@ void ggml_compute_forward_add( case GGML_TYPE_IQ3_XXS: case GGML_TYPE_IQ1_S: case GGML_TYPE_IQ1_M: + case GGML_TYPE_IQ2_NL: + case GGML_TYPE_IQ3_NL: case GGML_TYPE_IQ4_NL: case GGML_TYPE_IQ4_XS: case GGML_TYPE_IQ3_S: @@ -1135,6 +1137,8 @@ void ggml_compute_forward_add1( case GGML_TYPE_IQ3_XXS: case GGML_TYPE_IQ1_S: case GGML_TYPE_IQ1_M: + case GGML_TYPE_IQ2_NL: + case GGML_TYPE_IQ3_NL: case GGML_TYPE_IQ4_NL: case GGML_TYPE_IQ4_XS: case GGML_TYPE_IQ3_S: @@ -1265,6 +1269,8 @@ void ggml_compute_forward_acc( case GGML_TYPE_IQ3_XXS: case GGML_TYPE_IQ1_S: case GGML_TYPE_IQ1_M: + case GGML_TYPE_IQ2_NL: + case GGML_TYPE_IQ3_NL: case GGML_TYPE_IQ4_NL: case GGML_TYPE_IQ4_XS: case GGML_TYPE_IQ3_S: @@ -4461,6 +4467,8 @@ void ggml_compute_forward_out_prod( case GGML_TYPE_IQ3_XXS: case GGML_TYPE_IQ1_S: case GGML_TYPE_IQ1_M: + case GGML_TYPE_IQ2_NL: + case GGML_TYPE_IQ3_NL: case GGML_TYPE_IQ4_NL: case GGML_TYPE_IQ4_XS: case GGML_TYPE_IQ3_S: @@ -4738,6 +4746,8 @@ void ggml_compute_forward_set( case GGML_TYPE_IQ3_XXS: case GGML_TYPE_IQ1_S: case GGML_TYPE_IQ1_M: + case GGML_TYPE_IQ2_NL: + case GGML_TYPE_IQ3_NL: case GGML_TYPE_IQ4_NL: case GGML_TYPE_IQ4_XS: case GGML_TYPE_IQ3_S: @@ -4962,6 +4972,8 @@ void ggml_compute_forward_get_rows( case GGML_TYPE_IQ3_XXS: case GGML_TYPE_IQ1_S: case GGML_TYPE_IQ1_M: + case GGML_TYPE_IQ2_NL: + case GGML_TYPE_IQ3_NL: case GGML_TYPE_IQ4_NL: case GGML_TYPE_IQ4_XS: case GGML_TYPE_IQ3_S: From 92f87fda6bf1820a37b06e0a4c2ec90aafaa018c Mon Sep 17 00:00:00 2001 From: Ed Addario Date: Tue, 30 Jun 2026 23:04:41 +0100 Subject: [PATCH 10/21] Add new types to tolerance tests --- tests/test-quantize-fns.cpp | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/tests/test-quantize-fns.cpp b/tests/test-quantize-fns.cpp index 74978e446a3..b99204f3ab6 100644 --- a/tests/test-quantize-fns.cpp +++ b/tests/test-quantize-fns.cpp @@ -24,6 +24,7 @@ constexpr float MAX_QUANTIZATION_TOTAL_ERROR_3BITS_XXS = 0.0050f; constexpr float MAX_QUANTIZATION_TOTAL_ERROR_FP4 = 0.0030f; constexpr float MAX_DOT_PRODUCT_ERROR = 0.02f; constexpr float MAX_DOT_PRODUCT_ERROR_LOWBIT = 0.04f; +constexpr float MAX_DOT_PRODUCT_ERROR_NL2 = 0.06f; constexpr float MAX_DOT_PRODUCT_ERROR_FP4 = 0.03f; constexpr float MAX_DOT_PRODUCT_ERROR_BINARY = 0.40f; constexpr float MAX_DOT_PRODUCT_ERROR_TERNARY = 0.15f; @@ -160,8 +161,10 @@ static int test_vec_dot_q(bool verbose) { type == GGML_TYPE_TQ2_0 ? MAX_QUANTIZATION_TOTAL_ERROR_TERNARY : type == GGML_TYPE_Q2_K ? MAX_QUANTIZATION_TOTAL_ERROR_2BITS : type == GGML_TYPE_IQ2_S ? MAX_QUANTIZATION_TOTAL_ERROR_2BITS : + type == GGML_TYPE_IQ2_NL ? MAX_QUANTIZATION_TOTAL_ERROR_2BITS : type == GGML_TYPE_Q3_K ? MAX_QUANTIZATION_TOTAL_ERROR_3BITS : type == GGML_TYPE_IQ3_S ? MAX_QUANTIZATION_TOTAL_ERROR_3BITS : + type == GGML_TYPE_IQ3_NL ? MAX_QUANTIZATION_TOTAL_ERROR_3BITS : type == GGML_TYPE_IQ3_XXS ? MAX_QUANTIZATION_TOTAL_ERROR_3BITS_XXS : type == GGML_TYPE_NVFP4 ? MAX_QUANTIZATION_TOTAL_ERROR_FP4 : MAX_QUANTIZATION_TOTAL_ERROR; bool failed = !(total_error < max_quantization_error); @@ -178,8 +181,11 @@ static int test_vec_dot_q(bool verbose) { } const float vec_dot_error = dot_product_error(qfns, qfns_cpu, test_size, test_data.data(), test_data2.data()); - const float max_allowed_error = type == GGML_TYPE_Q2_K || type == GGML_TYPE_IQ2_XS || type == GGML_TYPE_IQ2_XXS || - type == GGML_TYPE_IQ3_XXS || type == GGML_TYPE_IQ3_S || type == GGML_TYPE_IQ2_S + const float max_allowed_error = type == GGML_TYPE_IQ2_NL + ? MAX_DOT_PRODUCT_ERROR_NL2 + : type == GGML_TYPE_Q2_K || type == GGML_TYPE_IQ2_XS || type == GGML_TYPE_IQ2_XXS || + type == GGML_TYPE_IQ3_XXS || type == GGML_TYPE_IQ3_S || type == GGML_TYPE_IQ2_S || + type == GGML_TYPE_IQ3_NL ? MAX_DOT_PRODUCT_ERROR_LOWBIT : type == GGML_TYPE_Q1_0 ? MAX_DOT_PRODUCT_ERROR_BINARY From 937be5a62e3cb419bcada432a43343264e83a3bc Mon Sep 17 00:00:00 2001 From: Ed Addario Date: Tue, 30 Jun 2026 23:07:01 +0100 Subject: [PATCH 11/21] Require imatrix for new types --- src/llama-quant.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/src/llama-quant.cpp b/src/llama-quant.cpp index 847e79f4655..830e7727bff 100644 --- a/src/llama-quant.cpp +++ b/src/llama-quant.cpp @@ -770,9 +770,11 @@ static bool tensor_requires_imatrix(const char * tensor_name, const ggml_type ds } switch (dst_type) { case GGML_TYPE_IQ3_XXS: + case GGML_TYPE_IQ3_NL: case GGML_TYPE_IQ2_XXS: case GGML_TYPE_IQ2_XS: case GGML_TYPE_IQ2_S: + case GGML_TYPE_IQ2_NL: case GGML_TYPE_IQ1_M: case GGML_TYPE_IQ1_S: return true; From dceb914b6b4e776d36faaddf7e957557566a9a48 Mon Sep 17 00:00:00 2001 From: Ed Addario Date: Wed, 1 Jul 2026 11:58:08 +0100 Subject: [PATCH 12/21] Add backend fallback warning --- ggml/src/ggml-metal/ggml-metal-device.m | 10 ++++++++-- src/llama-model-loader.cpp | 10 ++++++++++ 2 files changed, 18 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-metal/ggml-metal-device.m b/ggml/src/ggml-metal/ggml-metal-device.m index a7cbc60ebe4..c2986c7d67a 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.m +++ b/ggml/src/ggml-metal/ggml-metal-device.m @@ -490,6 +490,7 @@ void ggml_metal_encoder_debug_group_pop (ggml_metal_encoder_t encoder) { } void ggml_metal_encoder_set_pipeline(ggml_metal_encoder_t encoder, struct ggml_metal_pipeline_with_params pipeline) { + GGML_ASSERT(pipeline.pipeline && "missing Metal pipeline (kernel not found in library)"); [encoder->obj setComputePipelineState:pipeline.pipeline->obj]; } @@ -1266,7 +1267,10 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te case GGML_OP_SOLVE_TRI: case GGML_OP_MUL_MAT: case GGML_OP_MUL_MAT_ID: - return has_simdgroup_reduction && op->src[0]->type != GGML_TYPE_NVFP4; + return has_simdgroup_reduction && + op->src[0]->type != GGML_TYPE_NVFP4 && + op->src[0]->type != GGML_TYPE_IQ2_NL && + op->src[0]->type != GGML_TYPE_IQ3_NL; case GGML_OP_SET: case GGML_OP_CPY: case GGML_OP_DUP: @@ -1326,7 +1330,9 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te }; } case GGML_OP_GET_ROWS: - return op->src[0]->type != GGML_TYPE_NVFP4; + return op->src[0]->type != GGML_TYPE_NVFP4 && + op->src[0]->type != GGML_TYPE_IQ2_NL && + op->src[0]->type != GGML_TYPE_IQ3_NL; case GGML_OP_SET_ROWS: { if (op->src[0]->type != GGML_TYPE_F32) { diff --git a/src/llama-model-loader.cpp b/src/llama-model-loader.cpp index 474cabdfc09..4ceae14589e 100644 --- a/src/llama-model-loader.cpp +++ b/src/llama-model-loader.cpp @@ -1033,12 +1033,22 @@ static bool weight_buft_supported(const llama_hparams & hparams, ggml_tensor * w // find the first buffer type in the list that can use the tensor static ggml_backend_buffer_type_t select_weight_buft(const llama_hparams & hparams, ggml_tensor * tensor, ggml_op op, const buft_list_t * buft_list) { GGML_ASSERT(!buft_list->empty()); + bool skipped_gpu = false; for (const auto & cur : *buft_list) { ggml_backend_dev_t cur_dev = cur.first; ggml_backend_buffer_type_t cur_buft = cur.second; if (weight_buft_supported(hparams, tensor, op, cur_buft, cur_dev)) { + if (skipped_gpu) { + LLAMA_LOG_WARN("%s: no GPU backend supports operation %s for type %s in tensor %s; offloading to CPU (reduced performance)\n", + __func__, ggml_op_name(op), ggml_type_name(tensor->type), tensor->name); + } + return cur_buft; } + + if (ggml_backend_dev_type(cur_dev) == GGML_BACKEND_DEVICE_TYPE_GPU) { + skipped_gpu = true; + } } return nullptr; From 3362ab3d0f8941f1794dbb7b68168304599555d7 Mon Sep 17 00:00:00 2001 From: Ed Addario Date: Wed, 1 Jul 2026 12:14:01 +0100 Subject: [PATCH 13/21] Exclude iq2_nl/iq3_nl from sycl backend --- ggml/src/ggml-sycl/ggml-sycl.cpp | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 41449db665e..2d5c61b687b 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -5455,6 +5455,15 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons struct ggml_tensor * a = op->src[0]; struct ggml_tensor * b = op->src[1]; + // No SYCL kernels yet for iq2_nl and iq3_nl; fallback to CPU + if (a->type == GGML_TYPE_IQ2_NL || a->type == GGML_TYPE_IQ3_NL) { + return false; + } + + if (a->ne[3] != b->ne[3]) { + return false; + } + if (a->ne[3] != b->ne[3]) { return false; } From 02b8e95b2b0688024fd89d85ea64cc9068a3a508 Mon Sep 17 00:00:00 2001 From: Ed Addario Date: Wed, 1 Jul 2026 16:51:36 +0100 Subject: [PATCH 14/21] Wire new types --- gguf-py/gguf/constants.py | 2 ++ include/llama.h | 2 ++ src/llama-model-loader.cpp | 4 ++++ src/llama-quant.cpp | 25 ++++++++++++++++++------- tools/quantize/quantize.cpp | 2 ++ 5 files changed, 28 insertions(+), 7 deletions(-) diff --git a/gguf-py/gguf/constants.py b/gguf-py/gguf/constants.py index 10e6fef5841..8a915600df9 100644 --- a/gguf-py/gguf/constants.py +++ b/gguf-py/gguf/constants.py @@ -4471,6 +4471,8 @@ class LlamaFileType(IntEnum): MOSTLY_MXFP4_MOE = 38 # except 1d tensors MOSTLY_NVFP4 = 39 # except 1d tensors MOSTLY_Q1_0 = 40 # except 1d tensors + MOSTLY_IQ2_NL = 41 # except 1d tensors + MOSTLY_IQ3_NL = 42 # except 1d tensors GUESSED = 1024 # not specified in the model file diff --git a/include/llama.h b/include/llama.h index f723c9f60cf..97ad81834f7 100644 --- a/include/llama.h +++ b/include/llama.h @@ -155,6 +155,8 @@ extern "C" { LLAMA_FTYPE_MOSTLY_MXFP4_MOE = 38, // except 1d tensors LLAMA_FTYPE_MOSTLY_NVFP4 = 39, // except 1d tensors LLAMA_FTYPE_MOSTLY_Q1_0 = 40, // except 1d tensors + LLAMA_FTYPE_MOSTLY_IQ2_NL = 41, // except 1d tensors + LLAMA_FTYPE_MOSTLY_IQ3_NL = 42, // except 1d tensors LLAMA_FTYPE_GUESSED = 1024, // not specified in the model file }; diff --git a/src/llama-model-loader.cpp b/src/llama-model-loader.cpp index 4ceae14589e..b7c5cfdb0e1 100644 --- a/src/llama-model-loader.cpp +++ b/src/llama-model-loader.cpp @@ -64,6 +64,8 @@ static std::string llama_model_ftype_name(llama_ftype ftype) { case LLAMA_FTYPE_MOSTLY_IQ3_XXS: return "IQ3_XXS - 3.0625 bpw"; 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_IQ2_NL: return "IQ2_NL - 2.5 bpw"; + case LLAMA_FTYPE_MOSTLY_IQ3_NL: return "IQ3_NL - 3.5 bpw"; case LLAMA_FTYPE_MOSTLY_IQ4_NL: return "IQ4_NL - 4.5 bpw"; case LLAMA_FTYPE_MOSTLY_IQ4_XS: return "IQ4_XS - 4.25 bpw"; case LLAMA_FTYPE_MOSTLY_IQ3_S: return "IQ3_S - 3.4375 bpw"; @@ -757,6 +759,8 @@ llama_model_loader::llama_model_loader( case GGML_TYPE_IQ3_XXS: ftype = LLAMA_FTYPE_MOSTLY_IQ3_XXS; break; case GGML_TYPE_IQ1_S: ftype = LLAMA_FTYPE_MOSTLY_IQ1_S; break; case GGML_TYPE_IQ1_M: ftype = LLAMA_FTYPE_MOSTLY_IQ1_M; break; + case GGML_TYPE_IQ2_NL: ftype = LLAMA_FTYPE_MOSTLY_IQ2_NL; break; + case GGML_TYPE_IQ3_NL: ftype = LLAMA_FTYPE_MOSTLY_IQ3_NL; break; case GGML_TYPE_IQ4_NL: ftype = LLAMA_FTYPE_MOSTLY_IQ4_NL; break; case GGML_TYPE_IQ4_XS: ftype = LLAMA_FTYPE_MOSTLY_IQ4_XS; break; case GGML_TYPE_IQ3_S: ftype = LLAMA_FTYPE_MOSTLY_IQ3_S; break; diff --git a/src/llama-quant.cpp b/src/llama-quant.cpp index 830e7727bff..d3a63a15d58 100644 --- a/src/llama-quant.cpp +++ b/src/llama-quant.cpp @@ -375,10 +375,10 @@ static ggml_type tensor_type_fallback(quantize_state_impl & qs, const ggml_tenso case GGML_TYPE_IQ1_S: case GGML_TYPE_IQ1_M: case GGML_TYPE_IQ2_XXS: - case GGML_TYPE_IQ2_XS: - case GGML_TYPE_IQ2_S: + case GGML_TYPE_IQ2_XS: // types on the right: block size 32 + case GGML_TYPE_IQ2_S: return_type = GGML_TYPE_IQ2_NL; break; case GGML_TYPE_IQ3_XXS: - case GGML_TYPE_IQ3_S: // types on the right: block size 32 + case GGML_TYPE_IQ3_S: return_type = GGML_TYPE_IQ3_NL; break; case GGML_TYPE_IQ4_XS: return_type = GGML_TYPE_IQ4_NL; break; case GGML_TYPE_Q2_K: case GGML_TYPE_Q3_K: @@ -451,7 +451,7 @@ static ggml_type llama_tensor_get_type_impl(quantize_state_impl & qs, ggml_type } else if (ftype == LLAMA_FTYPE_MOSTLY_IQ2_XXS || ftype == LLAMA_FTYPE_MOSTLY_IQ2_XS || ftype == LLAMA_FTYPE_MOSTLY_IQ3_XXS || ftype == LLAMA_FTYPE_MOSTLY_IQ1_S || ftype == LLAMA_FTYPE_MOSTLY_IQ2_S || ftype == LLAMA_FTYPE_MOSTLY_IQ2_M || - ftype == LLAMA_FTYPE_MOSTLY_IQ1_M) { + ftype == LLAMA_FTYPE_MOSTLY_IQ2_NL || ftype == LLAMA_FTYPE_MOSTLY_IQ3_NL || ftype == LLAMA_FTYPE_MOSTLY_IQ1_M) { new_type = GGML_TYPE_Q5_K; } else if (new_type != GGML_TYPE_Q8_0) { @@ -477,6 +477,9 @@ static ggml_type llama_tensor_get_type_impl(quantize_state_impl & qs, ggml_type else if (ftype == LLAMA_FTYPE_MOSTLY_IQ2_S || ftype == LLAMA_FTYPE_MOSTLY_IQ2_M) { new_type = GGML_TYPE_IQ3_S; } + else if (ftype == LLAMA_FTYPE_MOSTLY_IQ2_NL || ftype == LLAMA_FTYPE_MOSTLY_IQ3_NL) { + new_type = GGML_TYPE_IQ3_NL; + } else if (ftype == LLAMA_FTYPE_MOSTLY_IQ3_XXS) { new_type = GGML_TYPE_IQ3_S; } @@ -531,6 +534,9 @@ static ggml_type llama_tensor_get_type_impl(quantize_state_impl & qs, ggml_type else if ((ftype == LLAMA_FTYPE_MOSTLY_IQ4_NL || ftype == LLAMA_FTYPE_MOSTLY_IQ4_XS) && qs.model.hparams.n_gqa() >= 4) { new_type = GGML_TYPE_Q5_K; } + else if ((ftype == LLAMA_FTYPE_MOSTLY_IQ2_NL || ftype == LLAMA_FTYPE_MOSTLY_IQ3_NL) && qs.model.hparams.n_gqa() >= 4) { + new_type = GGML_TYPE_IQ4_NL; + } else if ((ftype == LLAMA_FTYPE_MOSTLY_Q4_K_M || ftype == LLAMA_FTYPE_MOSTLY_Q5_K_M) && use_more_bits(qs.i_attention_wv, qs.n_attention_wv)) new_type = GGML_TYPE_Q6_K; else if (ftype == LLAMA_FTYPE_MOSTLY_Q4_K_S && qs.i_attention_wv < 4) new_type = GGML_TYPE_Q5_K; @@ -614,9 +620,10 @@ static ggml_type llama_tensor_get_type_impl(quantize_state_impl & qs, ggml_type if (arch != LLM_ARCH_FALCON) { if (qs.model.hparams.n_expert == 8) { if (ftype == LLAMA_FTYPE_MOSTLY_Q2_K || ftype == LLAMA_FTYPE_MOSTLY_IQ3_XS || ftype == LLAMA_FTYPE_MOSTLY_IQ3_XXS || - ftype == LLAMA_FTYPE_MOSTLY_Q3_K_S || ftype == LLAMA_FTYPE_MOSTLY_Q3_K_M || ftype == LLAMA_FTYPE_MOSTLY_IQ4_NL || - 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_Q3_K_S || ftype == LLAMA_FTYPE_MOSTLY_Q3_K_M || ftype == LLAMA_FTYPE_MOSTLY_IQ4_NL || + 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_IQ2_NL || + ftype == LLAMA_FTYPE_MOSTLY_IQ3_NL) { new_type = GGML_TYPE_Q5_K; } } else { @@ -817,8 +824,10 @@ ggml_type llama_ftype_get_default_type(llama_ftype ftype) { case LLAMA_FTYPE_MOSTLY_Q5_K_S: case LLAMA_FTYPE_MOSTLY_Q5_K_M: return GGML_TYPE_Q5_K; case LLAMA_FTYPE_MOSTLY_Q6_K: return GGML_TYPE_Q6_K; + // Ternary quants case LLAMA_FTYPE_MOSTLY_TQ1_0: return GGML_TYPE_TQ1_0; case LLAMA_FTYPE_MOSTLY_TQ2_0: return GGML_TYPE_TQ2_0; + // Non-linear quants case LLAMA_FTYPE_MOSTLY_IQ2_XXS: return GGML_TYPE_IQ2_XXS; case LLAMA_FTYPE_MOSTLY_IQ2_XS: return GGML_TYPE_IQ2_XS; case LLAMA_FTYPE_MOSTLY_IQ2_S: return GGML_TYPE_IQ2_XS; @@ -826,6 +835,8 @@ ggml_type llama_ftype_get_default_type(llama_ftype ftype) { case LLAMA_FTYPE_MOSTLY_IQ3_XXS: return GGML_TYPE_IQ3_XXS; case LLAMA_FTYPE_MOSTLY_IQ1_S: return GGML_TYPE_IQ1_S; case LLAMA_FTYPE_MOSTLY_IQ1_M: return GGML_TYPE_IQ1_M; + case LLAMA_FTYPE_MOSTLY_IQ2_NL: return GGML_TYPE_IQ2_NL; + case LLAMA_FTYPE_MOSTLY_IQ3_NL: return GGML_TYPE_IQ3_NL; case LLAMA_FTYPE_MOSTLY_IQ4_NL: return GGML_TYPE_IQ4_NL; case LLAMA_FTYPE_MOSTLY_IQ4_XS: return GGML_TYPE_IQ4_XS; case LLAMA_FTYPE_MOSTLY_IQ3_S: diff --git a/tools/quantize/quantize.cpp b/tools/quantize/quantize.cpp index 840eefc2f5a..9a3f4d3b7bd 100644 --- a/tools/quantize/quantize.cpp +++ b/tools/quantize/quantize.cpp @@ -56,6 +56,8 @@ static const std::vector QUANT_OPTIONS = { { "Q3_K_S", LLAMA_FTYPE_MOSTLY_Q3_K_S, " 3.41G, +1.6321 ppl @ Llama-3-8B", }, { "Q3_K_M", LLAMA_FTYPE_MOSTLY_Q3_K_M, " 3.74G, +0.6569 ppl @ Llama-3-8B", }, { "Q3_K_L", LLAMA_FTYPE_MOSTLY_Q3_K_L, " 4.03G, +0.5562 ppl @ Llama-3-8B", }, + { "IQ2_NL", LLAMA_FTYPE_MOSTLY_IQ2_NL, " 2.50 bpw non-linear quantization", }, + { "IQ3_NL", LLAMA_FTYPE_MOSTLY_IQ3_NL, " 3.50 bpw non-linear quantization", }, { "IQ4_NL", LLAMA_FTYPE_MOSTLY_IQ4_NL, " 4.50 bpw non-linear quantization", }, { "IQ4_XS", LLAMA_FTYPE_MOSTLY_IQ4_XS, " 4.25 bpw non-linear quantization", }, { "Q4_K", LLAMA_FTYPE_MOSTLY_Q4_K_M, "alias for Q4_K_M", }, From ae848177d103b05254b3f1cecd9a8e636fcdfd9d Mon Sep 17 00:00:00 2001 From: Ed Addario Date: Wed, 1 Jul 2026 17:32:05 +0100 Subject: [PATCH 15/21] Minor refactor --- ggml/src/ggml-common.h | 14 ++++++------ ggml/src/ggml-cpu/ops.cpp | 4 ++-- ggml/src/ggml-quants.c | 45 ++++++++++++++++++++------------------- ggml/src/ggml-quants.h | 12 +++++------ 4 files changed, 38 insertions(+), 37 deletions(-) diff --git a/ggml/src/ggml-common.h b/ggml/src/ggml-common.h index 68f58a55e51..3cb274b381a 100644 --- a/ggml/src/ggml-common.h +++ b/ggml/src/ggml-common.h @@ -434,13 +434,6 @@ typedef union { } iq1m_scale_t; // Non-linear quants -#define QK4_NL 32 -typedef struct { - ggml_half d; - uint8_t qs[QK4_NL/2]; -} block_iq4_nl; -static_assert(sizeof(block_iq4_nl) == sizeof(ggml_half) + QK4_NL/2, "wrong iq4_nl block size/padding"); - #define QK2_NL 32 typedef struct { ggml_half d; @@ -456,6 +449,13 @@ typedef struct { } block_iq3_nl; static_assert(sizeof(block_iq3_nl) == sizeof(ggml_half) + QK3_NL/8 + QK3_NL/4, "wrong iq3_nl block size/padding"); +#define QK4_NL 32 +typedef struct { + ggml_half d; + uint8_t qs[QK4_NL/2]; +} block_iq4_nl; +static_assert(sizeof(block_iq4_nl) == sizeof(ggml_half) + QK4_NL/2, "wrong iq4_nl block size/padding"); + typedef struct { ggml_half d; uint16_t scales_h; diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index 916efa41f82..860de1c4e4b 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -5700,12 +5700,12 @@ void ggml_compute_forward_clamp( case GGML_TYPE_IQ3_XXS: case GGML_TYPE_IQ1_S: case GGML_TYPE_IQ1_M: + case GGML_TYPE_IQ2_NL: + case GGML_TYPE_IQ3_NL: case GGML_TYPE_IQ4_NL: case GGML_TYPE_IQ4_XS: case GGML_TYPE_IQ3_S: case GGML_TYPE_IQ2_S: - case GGML_TYPE_IQ2_NL: - case GGML_TYPE_IQ3_NL: case GGML_TYPE_Q8_K: case GGML_TYPE_I8: case GGML_TYPE_I16: diff --git a/ggml/src/ggml-quants.c b/ggml/src/ggml-quants.c index c3d649dafc7..a8a4bd7613c 100644 --- a/ggml/src/ggml-quants.c +++ b/ggml/src/ggml-quants.c @@ -2650,24 +2650,6 @@ void dequantize_row_iq1_m(const block_iq1_m * GGML_RESTRICT x, float * GGML_REST } } -void dequantize_row_iq4_nl(const block_iq4_nl * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { - assert(k % QK4_NL == 0); - const int64_t nb = k / QK4_NL; - - for (int i = 0; i < nb; i++) { - - const uint8_t * qs = x[i].qs; - - const float d = GGML_FP16_TO_FP32(x[i].d); - for (int j = 0; j < QK4_NL/2; ++j) { - y[j+ 0] = d * kvalues_iq4nl[qs[j] & 0xf]; - y[j+QK4_NL/2] = d * kvalues_iq4nl[qs[j] >> 4]; - } - y += QK4_NL; - qs += QK4_NL/2; - } -} - void dequantize_row_iq2_nl(const block_iq2_nl * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { assert(k % QK2_NL == 0); const int64_t nb = k / QK2_NL; @@ -2705,6 +2687,24 @@ void dequantize_row_iq3_nl(const block_iq3_nl * GGML_RESTRICT x, float * GGML_RE } } +void dequantize_row_iq4_nl(const block_iq4_nl * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + assert(k % QK4_NL == 0); + const int64_t nb = k / QK4_NL; + + for (int i = 0; i < nb; i++) { + + const uint8_t * qs = x[i].qs; + + const float d = GGML_FP16_TO_FP32(x[i].d); + for (int j = 0; j < QK4_NL/2; ++j) { + y[j+ 0] = d * kvalues_iq4nl[qs[j] & 0xf]; + y[j+QK4_NL/2] = d * kvalues_iq4nl[qs[j] >> 4]; + } + y += QK4_NL; + qs += QK4_NL/2; + } +} + void dequantize_row_iq4_xs(const block_iq4_xs * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { assert(k % QK_K == 0); const int64_t nb = k / QK_K; @@ -5118,6 +5118,7 @@ void quantize_row_iq4_xs_ref(const float * GGML_RESTRICT x, block_iq4_xs * GGML_ quantize_iq4_xs(x, y, 1, k, NULL); } +// =============================== 2.5 & 3.5 bpw size_t quantize_iq2_nl(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { GGML_ASSERT(n_per_row%QK2_NL == 0); int64_t nblock = n_per_row/QK2_NL; @@ -5696,10 +5697,6 @@ bool ggml_validate_row_data(enum ggml_type type, const void * data, size_t nbyte { VALIDATE_ROW_DATA_D_F16_IMPL(block_iq4_xs, data, nb); } break; - case GGML_TYPE_IQ4_NL: - { - VALIDATE_ROW_DATA_D_F16_IMPL(block_iq4_nl, data, nb); - } break; case GGML_TYPE_IQ2_NL: { VALIDATE_ROW_DATA_D_F16_IMPL(block_iq2_nl, data, nb); @@ -5708,6 +5705,10 @@ bool ggml_validate_row_data(enum ggml_type type, const void * data, size_t nbyte { VALIDATE_ROW_DATA_D_F16_IMPL(block_iq3_nl, data, nb); } break; + case GGML_TYPE_IQ4_NL: + { + VALIDATE_ROW_DATA_D_F16_IMPL(block_iq4_nl, data, nb); + } break; case GGML_TYPE_I8: case GGML_TYPE_I16: diff --git a/ggml/src/ggml-quants.h b/ggml/src/ggml-quants.h index 4889a6a53af..4cb97510fcf 100644 --- a/ggml/src/ggml-quants.h +++ b/ggml/src/ggml-quants.h @@ -36,10 +36,10 @@ GGML_API void quantize_row_tq1_0_ref(const float * GGML_RESTRICT x, block_tq1_0 GGML_API void quantize_row_tq2_0_ref(const float * GGML_RESTRICT x, block_tq2_0 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_iq3_xxs_ref(const float * GGML_RESTRICT x, block_iq3_xxs * GGML_RESTRICT y, int64_t k); -GGML_API void quantize_row_iq4_nl_ref (const float * GGML_RESTRICT x, block_iq4_nl * GGML_RESTRICT y, int64_t k); -GGML_API void quantize_row_iq4_xs_ref (const float * GGML_RESTRICT x, block_iq4_xs * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_iq2_nl_ref (const float * GGML_RESTRICT x, block_iq2_nl * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_iq3_nl_ref (const float * GGML_RESTRICT x, block_iq3_nl * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_iq4_nl_ref (const float * GGML_RESTRICT x, block_iq4_nl * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_iq4_xs_ref (const float * GGML_RESTRICT x, block_iq4_xs * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_iq3_s_ref (const float * GGML_RESTRICT x, block_iq3_s * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_iq2_s_ref (const float * GGML_RESTRICT x, block_iq2_s * GGML_RESTRICT y, int64_t k); @@ -71,10 +71,10 @@ GGML_API void dequantize_row_iq2_s (const block_iq2_s * GGML_RESTRICT x, floa GGML_API void dequantize_row_iq3_xxs(const block_iq3_xxs * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_iq1_s (const block_iq1_s * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_iq1_m (const block_iq1_m * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); -GGML_API void dequantize_row_iq4_nl (const block_iq4_nl * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); -GGML_API void dequantize_row_iq4_xs (const block_iq4_xs * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_iq2_nl (const block_iq2_nl * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_iq3_nl (const block_iq3_nl * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_iq4_nl (const block_iq4_nl * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_iq4_xs (const block_iq4_xs * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_iq3_s (const block_iq3_s * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); // Quantization utilizing an importance matrix (a.k.a. "Activation aWare Quantization") @@ -84,10 +84,10 @@ GGML_API size_t quantize_iq2_s (const float * GGML_RESTRICT src, void * GGML_RE GGML_API size_t quantize_iq3_xxs(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_iq1_s (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_iq1_m (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); -GGML_API size_t quantize_iq4_nl (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); -GGML_API size_t quantize_iq4_xs (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_iq2_nl (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_iq3_nl (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); +GGML_API size_t quantize_iq4_nl (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); +GGML_API size_t quantize_iq4_xs (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_iq3_s (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); GGML_API size_t quantize_tq1_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); From 62f7a40247ef68fd9d8ccb604923d69dd33fdc50 Mon Sep 17 00:00:00 2001 From: Ed Addario Date: Wed, 1 Jul 2026 19:19:21 +0100 Subject: [PATCH 16/21] Enable testing new types --- tests/test-backend-ops.cpp | 33 +++++++++++++++++++++++++++------ 1 file changed, 27 insertions(+), 6 deletions(-) diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 15b50209c85..5d64db3b182 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -2396,13 +2396,23 @@ struct test_set_rows : public test_case { } double max_nmse_err() override { - if (type == GGML_TYPE_Q4_0 || type == GGML_TYPE_Q4_1 || type == GGML_TYPE_IQ4_NL || - type == GGML_TYPE_Q5_0 || type == GGML_TYPE_Q5_1 || type == GGML_TYPE_Q8_0) { + if (type == GGML_TYPE_IQ2_NL || type == GGML_TYPE_IQ3_NL || + type == GGML_TYPE_Q4_0 || type == GGML_TYPE_Q4_1 || type == GGML_TYPE_IQ4_NL || + type == GGML_TYPE_Q5_0 || type == GGML_TYPE_Q5_1 || type == GGML_TYPE_Q8_0) { // estimate what the max nmse error would be if one quantized value is // off by one. The test values are distributed in [-1,1], so it'll be // roughly (2.0 / 2^bits)^2, divided by the mean square value of the reference, // which is roughly 0.25 times the number of elements. double err_estimate = 1.0f/8.0f; + if (type == GGML_TYPE_IQ2_NL) { + // Using iq4_nl as reference, iq2_nl has only 4 non-linear levels (coarsest quant) + // so we increase error estimate to keep the same ratio (4 non-linear levels * 4.0f) + err_estimate *= 4.0f; + } + if (type == GGML_TYPE_IQ3_NL) { + // iq3_nl has 8 non-linear levels: coarser than the 16-level iq4_nl (8 non-linear levels * 2.0f) + err_estimate *= 2.0f; + } if (type == GGML_TYPE_Q5_0 || type == GGML_TYPE_Q5_1) { err_estimate /= 2.0f; } @@ -2915,13 +2925,22 @@ struct test_cpy : public test_case { if (type_src == type_dst) { return 0.0; } - if (type_dst == GGML_TYPE_Q4_0 || type_dst == GGML_TYPE_Q4_1 || type_dst == GGML_TYPE_IQ4_NL || - type_dst == GGML_TYPE_Q5_0 || type_dst == GGML_TYPE_Q5_1 || type_dst == GGML_TYPE_Q8_0) { + if (type_dst == GGML_TYPE_IQ2_NL || type_dst == GGML_TYPE_IQ3_NL || + type_dst == GGML_TYPE_Q4_0 || type_dst == GGML_TYPE_Q4_1 || type_dst == GGML_TYPE_IQ4_NL || + type_dst == GGML_TYPE_Q5_0 || type_dst == GGML_TYPE_Q5_1 || type_dst == GGML_TYPE_Q8_0) { // estimate what the max nmse error would be if one quantized value is // off by one. The test values are distributed in [-150,150], so it'll be // roughly (150*2.0 / 2^bits)^2, divided by the mean square value of the reference, // which is roughly 0.25*150^2 times the number of elements. double err_estimate = 1.0f/8.0f * 150.0f; + if (type_dst == GGML_TYPE_IQ2_NL) { + // 4 non-linear levels (coarsest) and more spread out so we increase error tolerance further + err_estimate *= 8.0f; + } + if (type_dst == GGML_TYPE_IQ3_NL) { + // 8 non-linear levels (coarser than iq4_nl's 16) and a bit more spread out so we increase error tolerance + err_estimate *= 4.0f; + } if (type_dst == GGML_TYPE_IQ4_NL) { // iq4_nl values are a bit more spread out err_estimate *= 2.0f; @@ -7639,7 +7658,8 @@ static const ggml_type all_types[] = { // GGML_TYPE_TQ1_0, GGML_TYPE_TQ2_0, // TODO: implement for all backends GGML_TYPE_IQ2_XXS, GGML_TYPE_IQ2_XS, GGML_TYPE_IQ2_S, GGML_TYPE_IQ3_XXS, GGML_TYPE_IQ1_S, GGML_TYPE_IQ1_M, - GGML_TYPE_IQ4_NL, GGML_TYPE_IQ3_S, GGML_TYPE_IQ4_XS, + GGML_TYPE_IQ2_NL, GGML_TYPE_IQ3_NL, GGML_TYPE_IQ4_NL, + GGML_TYPE_IQ3_S, GGML_TYPE_IQ4_XS, }; static const ggml_type base_types[] = { @@ -7664,7 +7684,8 @@ static const ggml_type other_types[] = { // GGML_TYPE_TQ1_0, GGML_TYPE_TQ2_0, // TODO: implement for all backends GGML_TYPE_IQ2_XS, GGML_TYPE_IQ2_S, GGML_TYPE_IQ3_XXS, GGML_TYPE_IQ1_S, GGML_TYPE_IQ1_M, - GGML_TYPE_IQ4_NL, GGML_TYPE_IQ3_S, GGML_TYPE_IQ4_XS, + GGML_TYPE_IQ2_NL, GGML_TYPE_IQ3_NL, GGML_TYPE_IQ4_NL, + GGML_TYPE_IQ3_S, GGML_TYPE_IQ4_XS, GGML_TYPE_BF16, }; From ac9a9f541afe287422c84728b8f1540d0cdbc07d Mon Sep 17 00:00:00 2001 From: Ed Addario Date: Sun, 19 Jul 2026 15:54:55 +0100 Subject: [PATCH 17/21] Finetune iq2_nl & iq3_nl codebooks --- ggml/src/ggml-common.h | 14 +++++++------- gguf-py/gguf/quants.py | 4 ++-- 2 files changed, 9 insertions(+), 9 deletions(-) diff --git a/ggml/src/ggml-common.h b/ggml/src/ggml-common.h index eb27de107fe..b2724f928c0 100644 --- a/ggml/src/ggml-common.h +++ b/ggml/src/ggml-common.h @@ -1131,17 +1131,17 @@ GGML_TABLE_BEGIN(uint32_t, iq3s_grid, 512) 0x0f090307, 0x0f090501, 0x0f090b01, 0x0f0b0505, 0x0f0b0905, 0x0f0d0105, 0x0f0d0703, 0x0f0f0101, GGML_TABLE_END() -// TODO: fix name to kvalues_iq4_nl -GGML_TABLE_BEGIN(int8_t, kvalues_iq4nl, 16) - -127, -104, -83, -65, -49, -35, -22, -10, 1, 13, 25, 38, 53, 69, 89, 113, -GGML_TABLE_END() - GGML_TABLE_BEGIN(int8_t, kvalues_iq2nl, 4) - -127, -38, 38, 127, + -127, -49, 9, 88, GGML_TABLE_END() GGML_TABLE_BEGIN(int8_t, kvalues_iq3nl, 8) - -127, -79, -45, -14, 14, 45, 79, 127, + -127, -78, -45, -18, 3, 30, 62, 107, +GGML_TABLE_END() + +// TODO: fix name to kvalues_iq4_nl +GGML_TABLE_BEGIN(int8_t, kvalues_iq4nl, 16) + -127, -104, -83, -65, -49, -35, -22, -10, 1, 13, 25, 38, 53, 69, 89, 113, GGML_TABLE_END() // e2m1 values (doubled), shared by MXFP4 and NVFP4 diff --git a/gguf-py/gguf/quants.py b/gguf-py/gguf/quants.py index dd4ca627836..cdf278fdefd 100644 --- a/gguf-py/gguf/quants.py +++ b/gguf-py/gguf/quants.py @@ -1329,7 +1329,7 @@ def dequantize_blocks(cls, blocks: np.ndarray) -> np.ndarray: class IQ2_NL(__Quant, qtype=GGMLQuantizationType.IQ2_NL): - kvalues = (-127, -38, 38, 127) + kvalues = (-127, -49, 9, 88) @classmethod def dequantize_blocks(cls, blocks: np.ndarray) -> np.ndarray: @@ -1349,7 +1349,7 @@ def dequantize_blocks(cls, blocks: np.ndarray) -> np.ndarray: class IQ3_NL(__Quant, qtype=GGMLQuantizationType.IQ3_NL): - kvalues = (-127, -79, -45, -14, 14, 45, 79, 127) + kvalues = (-127, -78, -45, -18, 3, 30, 62, 107) @classmethod def dequantize_blocks(cls, blocks: np.ndarray) -> np.ndarray: From 2cd66fb6e3c05313aa6c0d17080a43cf06717f6e Mon Sep 17 00:00:00 2001 From: Ed Addario Date: Sun, 16 Aug 2026 11:27:50 +0100 Subject: [PATCH 18/21] Update iq2_nl & iq3_nl ftype allocation --- src/llama-quant.cpp | 80 ++++++++++++++++++++++++++------------------- 1 file changed, 47 insertions(+), 33 deletions(-) diff --git a/src/llama-quant.cpp b/src/llama-quant.cpp index 062ccc25a46..29c62168a18 100644 --- a/src/llama-quant.cpp +++ b/src/llama-quant.cpp @@ -462,9 +462,15 @@ static ggml_type llama_tensor_get_type_impl(quantize_state_impl & qs, ggml_type else if (arch == LLM_ARCH_FALCON || nx % qk_k != 0) { new_type = GGML_TYPE_Q8_0; } + else if (ftype == LLAMA_FTYPE_MOSTLY_IQ2_NL) { + new_type = GGML_TYPE_Q5_0; + } + else if (ftype == LLAMA_FTYPE_MOSTLY_IQ3_NL) { + new_type = GGML_TYPE_Q5_1; + } else if (ftype == LLAMA_FTYPE_MOSTLY_IQ2_XXS || ftype == LLAMA_FTYPE_MOSTLY_IQ2_XS || ftype == LLAMA_FTYPE_MOSTLY_IQ3_XXS || ftype == LLAMA_FTYPE_MOSTLY_IQ1_S || ftype == LLAMA_FTYPE_MOSTLY_IQ2_S || ftype == LLAMA_FTYPE_MOSTLY_IQ2_M || - ftype == LLAMA_FTYPE_MOSTLY_IQ2_NL || ftype == LLAMA_FTYPE_MOSTLY_IQ3_NL || ftype == LLAMA_FTYPE_MOSTLY_IQ1_M) { + ftype == LLAMA_FTYPE_MOSTLY_IQ1_M) { new_type = GGML_TYPE_Q5_K; } else if (new_type != GGML_TYPE_Q8_0) { @@ -503,8 +509,8 @@ static ggml_type llama_tensor_get_type_impl(quantize_state_impl & qs, ggml_type } else if (ftype == LLAMA_FTYPE_MOSTLY_IQ2_XXS || ftype == LLAMA_FTYPE_MOSTLY_IQ2_XS || ftype == LLAMA_FTYPE_MOSTLY_IQ1_S || ftype == LLAMA_FTYPE_MOSTLY_IQ2_S || ftype == LLAMA_FTYPE_MOSTLY_IQ2_M || ftype == LLAMA_FTYPE_MOSTLY_IQ1_M) { if (category_is_attn_v(category)) { - if (qs.model.hparams.n_gqa() >= 4 || qs.model.hparams.n_expert >= 4) new_type = GGML_TYPE_Q4_K; - else new_type = ftype == LLAMA_FTYPE_MOSTLY_IQ2_S || ftype == LLAMA_FTYPE_MOSTLY_IQ2_M ? GGML_TYPE_IQ3_S : GGML_TYPE_Q2_K; + if (qs.model.hparams.n_gqa() >= 4 || qs.model.hparams.n_expert >= 4) { new_type = GGML_TYPE_Q4_K; } + else { new_type = ftype == LLAMA_FTYPE_MOSTLY_IQ2_S || ftype == LLAMA_FTYPE_MOSTLY_IQ2_M ? GGML_TYPE_IQ3_S : GGML_TYPE_Q2_K; } ++qs.i_attention_wv; } else if (qs.model.hparams.n_expert == 8 && category == tensor_category::ATTENTION_K) { @@ -520,8 +526,8 @@ static ggml_type llama_tensor_get_type_impl(quantize_state_impl & qs, ggml_type if (qs.model.hparams.n_expert == 8) { new_type = GGML_TYPE_Q5_K; } else { - if (ftype == LLAMA_FTYPE_MOSTLY_IQ1_S || ftype == LLAMA_FTYPE_MOSTLY_IQ1_M) new_type = GGML_TYPE_IQ2_XXS; - else if (ftype == LLAMA_FTYPE_MOSTLY_IQ2_S || ftype == LLAMA_FTYPE_MOSTLY_IQ2_M) new_type = GGML_TYPE_IQ3_S; + if (ftype == LLAMA_FTYPE_MOSTLY_IQ1_S || ftype == LLAMA_FTYPE_MOSTLY_IQ1_M) { new_type = GGML_TYPE_IQ2_XXS; } + else if (ftype == LLAMA_FTYPE_MOSTLY_IQ2_S || ftype == LLAMA_FTYPE_MOSTLY_IQ2_M) { new_type = GGML_TYPE_IQ3_S; } } } } else if (category_is_attn_v(category)) { @@ -543,33 +549,34 @@ static ggml_type llama_tensor_get_type_impl(quantize_state_impl & qs, ggml_type else if (ftype == LLAMA_FTYPE_MOSTLY_Q3_K_M) { new_type = qs.i_attention_wv < 2 ? GGML_TYPE_Q5_K : GGML_TYPE_Q4_K; } - else if (ftype == LLAMA_FTYPE_MOSTLY_Q3_K_L) new_type = GGML_TYPE_Q5_K; + 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) && qs.model.hparams.n_gqa() >= 4) { new_type = GGML_TYPE_Q5_K; } - else if ((ftype == LLAMA_FTYPE_MOSTLY_IQ2_NL || ftype == LLAMA_FTYPE_MOSTLY_IQ3_NL) && qs.model.hparams.n_gqa() >= 4) { + else if (ftype == LLAMA_FTYPE_MOSTLY_IQ3_NL) { new_type = GGML_TYPE_IQ4_NL; } + else if (ftype == LLAMA_FTYPE_MOSTLY_IQ2_NL) { + new_type = qs.model.hparams.n_gqa() >= 4 || qs.model.hparams.n_expert >= 4 ? GGML_TYPE_IQ4_NL : GGML_TYPE_IQ3_NL; + } else if ((ftype == LLAMA_FTYPE_MOSTLY_Q4_K_M || ftype == LLAMA_FTYPE_MOSTLY_Q5_K_M) && - use_more_bits(qs.i_attention_wv, qs.n_attention_wv)) new_type = GGML_TYPE_Q6_K; - else if (ftype == LLAMA_FTYPE_MOSTLY_Q4_K_S && qs.i_attention_wv < 4) new_type = GGML_TYPE_Q5_K; + use_more_bits(qs.i_attention_wv, qs.n_attention_wv)) { new_type = GGML_TYPE_Q6_K; } + else if (ftype == LLAMA_FTYPE_MOSTLY_Q4_K_S && qs.i_attention_wv < 4) { new_type = GGML_TYPE_Q5_K; } if (qs.model.type == LLM_TYPE_70B) { // In the 70B model we have 8 heads sharing the same attn_v weights. As a result, the attn_v.weight tensor is // 8x smaller compared to attn_q.weight. Hence, we can get a nice boost in quantization accuracy with // nearly negligible increase in model size by quantizing this tensor with more bits: - if (new_type == GGML_TYPE_Q3_K || new_type == GGML_TYPE_Q4_K) new_type = GGML_TYPE_Q5_K; + if (new_type == GGML_TYPE_Q3_K || new_type == GGML_TYPE_Q4_K) { new_type = GGML_TYPE_Q5_K; } } if (qs.model.hparams.n_expert == 8) { // for the 8-expert model, bumping this to Q8_0 trades just ~128MB - // TODO: explore better strategies - new_type = GGML_TYPE_Q8_0; + new_type = ftype == LLAMA_FTYPE_MOSTLY_IQ2_NL || ftype == LLAMA_FTYPE_MOSTLY_IQ3_NL ? GGML_TYPE_IQ4_NL : GGML_TYPE_Q8_0; } ++qs.i_attention_wv; } else if (category == tensor_category::ATTENTION_K) { if (qs.model.hparams.n_expert == 8) { // for the 8-expert model, bumping this to Q8_0 trades just ~128MB - // TODO: explore better strategies - new_type = GGML_TYPE_Q8_0; + new_type = ftype == LLAMA_FTYPE_MOSTLY_IQ2_NL || ftype == LLAMA_FTYPE_MOSTLY_IQ3_NL ? GGML_TYPE_IQ4_NL : GGML_TYPE_Q8_0; } else if (ftype == LLAMA_FTYPE_MOSTLY_IQ3_XS) { new_type = GGML_TYPE_IQ3_XXS; @@ -587,9 +594,9 @@ static ggml_type llama_tensor_get_type_impl(quantize_state_impl & qs, ggml_type } else if (category == tensor_category::FFN_DOWN) { auto info = layer_info(qs.i_ffn_down, qs.n_ffn_down, name.c_str()); int i_layer = info.first, n_layer = info.second; - if (ftype == LLAMA_FTYPE_MOSTLY_Q2_K) new_type = GGML_TYPE_Q3_K; + if (ftype == LLAMA_FTYPE_MOSTLY_Q2_K) { new_type = GGML_TYPE_Q3_K; } else if (ftype == LLAMA_FTYPE_MOSTLY_Q2_K_S) { - if (i_layer < n_layer/8) new_type = GGML_TYPE_Q4_K; + if (i_layer < n_layer/8) { new_type = GGML_TYPE_Q4_K; } } else if (ftype == LLAMA_FTYPE_MOSTLY_IQ3_XXS && !qs.has_imatrix) { new_type = i_layer < n_layer/8 ? GGML_TYPE_Q4_K : GGML_TYPE_Q3_K; @@ -611,13 +618,19 @@ static ggml_type llama_tensor_get_type_impl(quantize_state_impl & qs, ggml_type new_type = i_layer < n_layer/16 ? GGML_TYPE_Q6_K : use_more_bits(i_layer, n_layer) ? GGML_TYPE_Q5_K : GGML_TYPE_Q4_K; } else { - if (use_more_bits(i_layer, n_layer)) new_type = GGML_TYPE_Q6_K; + if (use_more_bits(i_layer, n_layer)) { new_type = GGML_TYPE_Q6_K; } } } else if (i_layer < n_layer/8 && (ftype == LLAMA_FTYPE_MOSTLY_IQ4_NL || ftype == LLAMA_FTYPE_MOSTLY_IQ4_XS) && !qs.has_imatrix) { 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; + else if (ftype == LLAMA_FTYPE_MOSTLY_IQ2_NL && i_layer < n_layer / 8) { + new_type = GGML_TYPE_IQ3_NL; + } + else if (ftype == LLAMA_FTYPE_MOSTLY_IQ3_NL && (i_layer < n_layer / 8 || (qs.model.hparams.n_expert == 8 && use_more_bits(i_layer, n_layer)))) { + new_type = GGML_TYPE_IQ4_NL; + } + else if (ftype == LLAMA_FTYPE_MOSTLY_Q5_K_M && use_more_bits(i_layer, n_layer)) { new_type = GGML_TYPE_Q6_K; } else if (ftype == LLAMA_FTYPE_MOSTLY_Q4_K_S && arch != LLM_ARCH_FALCON && i_layer < n_layer/8) { new_type = GGML_TYPE_Q5_K; } @@ -632,30 +645,31 @@ static ggml_type llama_tensor_get_type_impl(quantize_state_impl & qs, ggml_type } else if (category == tensor_category::ATTENTION_OUTPUT) { if (arch != LLM_ARCH_FALCON) { if (qs.model.hparams.n_expert == 8) { - if (ftype == LLAMA_FTYPE_MOSTLY_Q2_K || ftype == LLAMA_FTYPE_MOSTLY_IQ3_XS || ftype == LLAMA_FTYPE_MOSTLY_IQ3_XXS || - ftype == LLAMA_FTYPE_MOSTLY_Q3_K_S || ftype == LLAMA_FTYPE_MOSTLY_Q3_K_M || ftype == LLAMA_FTYPE_MOSTLY_IQ4_NL || - 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_IQ2_NL || - ftype == LLAMA_FTYPE_MOSTLY_IQ3_NL) { - new_type = GGML_TYPE_Q5_K; - } + if (ftype == LLAMA_FTYPE_MOSTLY_IQ2_NL) { new_type = GGML_TYPE_Q5_0; } + else if (ftype == LLAMA_FTYPE_MOSTLY_IQ3_NL) { new_type = GGML_TYPE_Q5_1; } + else if (ftype == LLAMA_FTYPE_MOSTLY_Q2_K || ftype == LLAMA_FTYPE_MOSTLY_IQ3_XS || ftype == LLAMA_FTYPE_MOSTLY_IQ3_XXS || + ftype == LLAMA_FTYPE_MOSTLY_Q3_K_S || ftype == LLAMA_FTYPE_MOSTLY_Q3_K_M || ftype == LLAMA_FTYPE_MOSTLY_IQ4_NL || + 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) { new_type = GGML_TYPE_Q5_K; } } else { - if (ftype == LLAMA_FTYPE_MOSTLY_Q2_K ) new_type = GGML_TYPE_Q3_K; - else if (ftype == LLAMA_FTYPE_MOSTLY_IQ3_XXS) new_type = GGML_TYPE_IQ3_S; - else if (ftype == LLAMA_FTYPE_MOSTLY_Q3_K_M ) new_type = GGML_TYPE_Q4_K; - else if (ftype == LLAMA_FTYPE_MOSTLY_Q3_K_L ) new_type = GGML_TYPE_Q5_K; - else if (ftype == LLAMA_FTYPE_MOSTLY_IQ3_M ) new_type = GGML_TYPE_Q4_K; + if (ftype == LLAMA_FTYPE_MOSTLY_Q2_K ) { new_type = GGML_TYPE_Q3_K; } + else if (ftype == LLAMA_FTYPE_MOSTLY_IQ3_XXS) { new_type = GGML_TYPE_IQ3_S; } + else if (ftype == LLAMA_FTYPE_MOSTLY_Q3_K_M ) { new_type = GGML_TYPE_Q4_K; } + else if (ftype == LLAMA_FTYPE_MOSTLY_Q3_K_L ) { new_type = GGML_TYPE_Q5_K; } + else if (ftype == LLAMA_FTYPE_MOSTLY_IQ3_M ) { new_type = GGML_TYPE_Q4_K; } + else if (ftype == LLAMA_FTYPE_MOSTLY_IQ2_NL ) { new_type = GGML_TYPE_IQ3_NL; } + else if (ftype == LLAMA_FTYPE_MOSTLY_IQ3_NL ) { new_type = GGML_TYPE_IQ4_NL; } } } else { - if (ftype == LLAMA_FTYPE_MOSTLY_Q3_K_L) new_type = GGML_TYPE_Q4_K; + if (ftype == LLAMA_FTYPE_MOSTLY_Q3_K_L) { new_type = GGML_TYPE_Q4_K; } } } else if (category == tensor_category::ATTENTION_QKV) { if (ftype == LLAMA_FTYPE_MOSTLY_Q3_K_M || ftype == LLAMA_FTYPE_MOSTLY_Q3_K_L || ftype == LLAMA_FTYPE_MOSTLY_IQ3_M) { new_type = GGML_TYPE_Q4_K; } - else if (ftype == LLAMA_FTYPE_MOSTLY_Q4_K_M) new_type = GGML_TYPE_Q5_K; - else if (ftype == LLAMA_FTYPE_MOSTLY_Q5_K_M) new_type = GGML_TYPE_Q6_K; + else if (ftype == LLAMA_FTYPE_MOSTLY_Q4_K_M) { new_type = GGML_TYPE_Q5_K; } + else if (ftype == LLAMA_FTYPE_MOSTLY_Q5_K_M) { new_type = GGML_TYPE_Q6_K; } } else if (category == tensor_category::FFN_GATE) { auto info = layer_info(qs.i_ffn_gate, qs.n_ffn_gate, name.c_str()); From b042012c3be524bfc8e5cb5a08e322cfb771f80f Mon Sep 17 00:00:00 2001 From: Ed Addario <29247825+EAddario@users.noreply.github.com> Date: Tue, 18 Aug 2026 15:07:28 +0100 Subject: [PATCH 19/21] Update ggml/src/ggml-sycl/ggml-sycl.cpp MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Co-authored-by: Sigbjørn Skjæret --- ggml/src/ggml-sycl/ggml-sycl.cpp | 4 ---- 1 file changed, 4 deletions(-) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 7a1f32f8b43..881b08fcc81 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -5973,10 +5973,6 @@ static bool do_ggml_backend_sycl_device_supports_op(ggml_backend_dev_t dev, cons return false; } - if (a->ne[3] != b->ne[3]) { - return false; - } - ggml_type src0_type = op->src[0]->type; // TODO: The configuration below needs more work to be supported with oneDNN From 93d328493ce0ef3ab924e36f3f1fc190da4a7e28 Mon Sep 17 00:00:00 2001 From: Ed Addario Date: Mon, 31 Aug 2026 23:26:11 +0200 Subject: [PATCH 20/21] Performance optimization --- ggml/src/ggml-cpu/arch-fallback.h | 5 -- ggml/src/ggml-cpu/arch/arm/quants.c | 123 ++++++++++++++++++++++++++++ ggml/src/ggml-cpu/arch/x86/quants.c | 110 +++++++++++++++++++++++++ ggml/src/ggml-quants.c | 19 +++-- 4 files changed, 244 insertions(+), 13 deletions(-) diff --git a/ggml/src/ggml-cpu/arch-fallback.h b/ggml/src/ggml-cpu/arch-fallback.h index 6624c2eb282..8917666ddba 100644 --- a/ggml/src/ggml-cpu/arch-fallback.h +++ b/ggml/src/ggml-cpu/arch-fallback.h @@ -74,9 +74,6 @@ #define ggml_gemm_q8_0_4x4_q8_0_generic ggml_gemm_q8_0_4x4_q8_0 #define ggml_gemm_q8_0_4x8_q8_0_generic ggml_gemm_q8_0_4x8_q8_0 #elif defined(__aarch64__) || defined(__arm__) || defined(_M_ARM) || defined(_M_ARM64) -// quants.c -#define ggml_vec_dot_iq2_nl_q8_0_generic ggml_vec_dot_iq2_nl_q8_0 -#define ggml_vec_dot_iq3_nl_q8_0_generic ggml_vec_dot_iq3_nl_q8_0 // repack.cpp #define ggml_quantize_mat_q8_K_4x4_generic ggml_quantize_mat_q8_K_4x4 #define ggml_quantize_mat_q8_K_4x8_generic ggml_quantize_mat_q8_K_4x8 @@ -89,8 +86,6 @@ #elif defined(__x86_64__) || defined(__i386__) || defined(_M_IX86) || defined(_M_X64) // quants.c #define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0 -#define ggml_vec_dot_iq2_nl_q8_0_generic ggml_vec_dot_iq2_nl_q8_0 -#define ggml_vec_dot_iq3_nl_q8_0_generic ggml_vec_dot_iq3_nl_q8_0 // repack.cpp #define ggml_quantize_mat_q8_0_4x4_generic ggml_quantize_mat_q8_0_4x4 #define ggml_quantize_mat_q8_K_4x4_generic ggml_quantize_mat_q8_K_4x4 diff --git a/ggml/src/ggml-cpu/arch/arm/quants.c b/ggml/src/ggml-cpu/arch/arm/quants.c index b988abf9963..4646bb233d5 100644 --- a/ggml/src/ggml-cpu/arch/arm/quants.c +++ b/ggml/src/ggml-cpu/arch/arm/quants.c @@ -4193,6 +4193,129 @@ void ggml_vec_dot_iq1_m_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const vo #endif } +void ggml_vec_dot_iq2_nl_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) { + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + assert(n % QK2_NL == 0); + static_assert(QK2_NL == QK8_0, "QK2_NL and QK8_0 must be the same"); + + const block_iq2_nl * GGML_RESTRICT x = vx; + const block_q8_0 * GGML_RESTRICT y = vy; + + const int nb = n / QK2_NL; + + int ib = 0; + float sumf = 0; + +#if defined __ARM_NEON + int32_t aux32; + memcpy(&aux32, kvalues_iq2nl, sizeof(aux32)); + + // splat the 4 entry LUT over the whole table, only indices 0..3 are used + const int8x16_t values = vreinterpretq_s8_s32(vdupq_n_s32(aux32)); + const uint8x16_t m3 = vdupq_n_u8(0x03); + + // a qs byte holds 4 weights that are 8 apart, so each group of 8 lanes needs its own shift + const int8x16_t sh_1 = vcombine_s8(vdup_n_s8( 0), vdup_n_s8(-2)); + const int8x16_t sh_2 = vcombine_s8(vdup_n_s8(-4), vdup_n_s8(-6)); + + for (; ib < nb; ++ib) { + const uint8x8_t q2bits = vld1_u8(x[ib].qs); + const uint8x16_t qs = vcombine_u8(q2bits, q2bits); + + const int8x16_t q2b_1 = ggml_vqtbl1q_s8(values, vandq_u8(vshlq_u8(qs, sh_1), m3)); + const int8x16_t q2b_2 = ggml_vqtbl1q_s8(values, vandq_u8(vshlq_u8(qs, sh_2), m3)); + + const int32x4_t prod = ggml_vdotq_s32(ggml_vdotq_s32(vdupq_n_s32(0), q2b_1, vld1q_s8(y[ib].qs)), + q2b_2, vld1q_s8(y[ib].qs + 16)); + + sumf += GGML_CPU_FP16_TO_FP32(x[ib].d) * GGML_CPU_FP16_TO_FP32(y[ib].d) * vaddvq_s32(prod); + } + +#endif + for (; ib < nb; ++ib) { + const float d = GGML_CPU_FP16_TO_FP32(y[ib].d)*GGML_CPU_FP16_TO_FP32(x[ib].d); + int sumi = 0; + for (int j = 0; j < QK2_NL/4; ++j) { + for (int g = 0; g < 4; ++g) { + sumi += y[ib].qs[j + g*(QK2_NL/4)] * kvalues_iq2nl[(x[ib].qs[j] >> 2*g) & 3]; + } + } + sumf += d * sumi; + } + *s = sumf; +} + +void ggml_vec_dot_iq3_nl_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) { + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + assert(n % QK3_NL == 0); + static_assert(QK3_NL == QK8_0, "QK3_NL and QK8_0 must be the same"); + + const block_iq3_nl * GGML_RESTRICT x = vx; + const block_q8_0 * GGML_RESTRICT y = vy; + + const int nb = n / QK3_NL; + + int ib = 0; + float sumf = 0; + +#if defined __ARM_NEON + const int8x8_t values8 = vld1_s8(kvalues_iq3nl); + const int8x16_t values = vcombine_s8(values8, values8); + + const uint8x16_t m3 = vdupq_n_u8(0x03); + const uint8x16_t m4 = vdupq_n_u8(0x04); + + const int8x16_t sh_1 = vcombine_s8(vdup_n_s8( 0), vdup_n_s8(-2)); + const int8x16_t sh_2 = vcombine_s8(vdup_n_s8(-4), vdup_n_s8(-6)); + + // qh byte g holds the high bit of the 8 weights of group g, one bit per lane + const uint8x16_t bit = vreinterpretq_u8_u64(vdupq_n_u64(0x8040201008040201ULL)); + const uint8x16_t hsel_1 = vcombine_u8(vdup_n_u8(0), vdup_n_u8(1)); + const uint8x16_t hsel_2 = vcombine_u8(vdup_n_u8(2), vdup_n_u8(3)); + + for (; ib < nb; ++ib) { + const uint8x8_t q3bits = vld1_u8(x[ib].qs); + const uint8x16_t qs = vcombine_u8(q3bits, q3bits); + + uint32_t aux32; + memcpy(&aux32, x[ib].qh, sizeof(aux32)); + const uint8x16_t qh = vreinterpretq_u8_u32(vdupq_n_u32(aux32)); + + const uint8x16_t idx_1 = vorrq_u8(vandq_u8(vshlq_u8(qs, sh_1), m3), + vandq_u8(vtstq_u8(ggml_vqtbl1q_u8(qh, hsel_1), bit), m4)); + const uint8x16_t idx_2 = vorrq_u8(vandq_u8(vshlq_u8(qs, sh_2), m3), + vandq_u8(vtstq_u8(ggml_vqtbl1q_u8(qh, hsel_2), bit), m4)); + + const int32x4_t prod = ggml_vdotq_s32(ggml_vdotq_s32(vdupq_n_s32(0), + ggml_vqtbl1q_s8(values, idx_1), vld1q_s8(y[ib].qs)), + ggml_vqtbl1q_s8(values, idx_2), vld1q_s8(y[ib].qs + 16)); + + sumf += GGML_CPU_FP16_TO_FP32(x[ib].d) * GGML_CPU_FP16_TO_FP32(y[ib].d) * vaddvq_s32(prod); + } + +#endif + for (; ib < nb; ++ib) { + const float d = GGML_CPU_FP16_TO_FP32(y[ib].d)*GGML_CPU_FP16_TO_FP32(x[ib].d); + int sumi = 0; + for (int j = 0; j < QK3_NL/4; ++j) { + for (int g = 0; g < 4; ++g) { + const int idx = ((x[ib].qs[j] >> 2*g) & 3) | (((x[ib].qh[g] >> j) & 1) << 2); + sumi += y[ib].qs[j + g*(QK3_NL/4)] * kvalues_iq3nl[idx]; + } + } + sumf += d * sumi; + } + *s = sumf; +} + void ggml_vec_dot_iq4_nl_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) { assert(nrc == 1); UNUSED(nrc); diff --git a/ggml/src/ggml-cpu/arch/x86/quants.c b/ggml/src/ggml-cpu/arch/x86/quants.c index ea54cfe44ce..08d6be07c76 100644 --- a/ggml/src/ggml-cpu/arch/x86/quants.c +++ b/ggml/src/ggml-cpu/arch/x86/quants.c @@ -85,6 +85,14 @@ static inline __m256i bytes_from_bits_32(const uint8_t * x) { return _mm256_cmpeq_epi8(bytes, _mm256_set1_epi64x(-1)); } +// Unpack 32 2-bit fields into 32 bytes, grouped 8 apart as the iq2_nl/iq3_nl qs packing stores them +static inline __m256i bytes_from_2bits_32(const uint8_t * x) { + uint64_t x64; + memcpy(&x64, x, sizeof(uint64_t)); + const __m256i shifted = _mm256_srlv_epi64(_mm256_set1_epi64x(x64), _mm256_set_epi64x(6, 4, 2, 0)); + return _mm256_and_si256(shifted, _mm256_set1_epi8(0x03)); +} + // Unpack 32 4-bit fields into 32 bytes // The output vector contains 32 bytes, each one in [ 0 .. 15 ] interval static inline __m256i bytes_from_nibbles_32(const uint8_t * rsi) @@ -3917,6 +3925,108 @@ void ggml_vec_dot_iq1_m_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const vo #endif } +void ggml_vec_dot_iq2_nl_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) { + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + assert(n % QK2_NL == 0); + static_assert(QK2_NL == QK8_0, "QK2_NL and QK8_0 must be the same"); + + const block_iq2_nl * GGML_RESTRICT x = vx; + const block_q8_0 * GGML_RESTRICT y = vy; + + const int nb = n / QK2_NL; + + int ib = 0; + float sumf = 0; + +#if defined __AVX2__ + int32_t aux32; + memcpy(&aux32, kvalues_iq2nl, sizeof(aux32)); + + // splat the 4 entry LUT over both 128 bit lanes, only indices 0..3 are used + const __m256i values = _mm256_set1_epi32(aux32); + const __m256i mone = _mm256_set1_epi16(1); + + __m256 accum = _mm256_setzero_ps(); + for (; ib < nb; ++ib) { + const __m256i q2b = _mm256_shuffle_epi8(values, bytes_from_2bits_32(x[ib].qs)); + const __m256i q8b = _mm256_loadu_si256((const __m256i *)y[ib].qs); + const __m256i p = _mm256_madd_epi16(mul_add_epi8(q2b, q8b), mone); + accum = _mm256_fmadd_ps(_mm256_set1_ps(GGML_CPU_FP16_TO_FP32(y[ib].d)*GGML_CPU_FP16_TO_FP32(x[ib].d)), + _mm256_cvtepi32_ps(p), accum); + } + + sumf = hsum_float_8(accum); + +#endif + for (; ib < nb; ++ib) { + const float d = GGML_CPU_FP16_TO_FP32(y[ib].d)*GGML_CPU_FP16_TO_FP32(x[ib].d); + int sumi = 0; + for (int j = 0; j < QK2_NL/4; ++j) { + for (int g = 0; g < 4; ++g) { + sumi += y[ib].qs[j + g*(QK2_NL/4)] * kvalues_iq2nl[(x[ib].qs[j] >> 2*g) & 3]; + } + } + sumf += d * sumi; + } + *s = sumf; +} + +void ggml_vec_dot_iq3_nl_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) { + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + assert(n % QK3_NL == 0); + static_assert(QK3_NL == QK8_0, "QK3_NL and QK8_0 must be the same"); + + const block_iq3_nl * GGML_RESTRICT x = vx; + const block_q8_0 * GGML_RESTRICT y = vy; + + const int nb = n / QK3_NL; + + int ib = 0; + float sumf = 0; + +#if defined __AVX2__ + const __m128i values128 = _mm_loadl_epi64((const __m128i *)kvalues_iq3nl); + const __m256i values = MM256_SET_M128I(values128, values128); + const __m256i m4 = _mm256_set1_epi8(0x04); + const __m256i mone = _mm256_set1_epi16(1); + + __m256 accum = _mm256_setzero_ps(); + for (; ib < nb; ++ib) { + // qh keeps one high bit per weight in weight order, which is what bytes_from_bits_32 expands to + const __m256i qh = _mm256_and_si256(bytes_from_bits_32(x[ib].qh), m4); + const __m256i idx = _mm256_or_si256(bytes_from_2bits_32(x[ib].qs), qh); + const __m256i q3b = _mm256_shuffle_epi8(values, idx); + const __m256i q8b = _mm256_loadu_si256((const __m256i *)y[ib].qs); + const __m256i p = _mm256_madd_epi16(mul_add_epi8(q3b, q8b), mone); + accum = _mm256_fmadd_ps(_mm256_set1_ps(GGML_CPU_FP16_TO_FP32(y[ib].d)*GGML_CPU_FP16_TO_FP32(x[ib].d)), + _mm256_cvtepi32_ps(p), accum); + } + + sumf = hsum_float_8(accum); + +#endif + for (; ib < nb; ++ib) { + const float d = GGML_CPU_FP16_TO_FP32(y[ib].d)*GGML_CPU_FP16_TO_FP32(x[ib].d); + int sumi = 0; + for (int j = 0; j < QK3_NL/4; ++j) { + for (int g = 0; g < 4; ++g) { + const int idx = ((x[ib].qs[j] >> 2*g) & 3) | (((x[ib].qh[g] >> j) & 1) << 2); + sumi += y[ib].qs[j + g*(QK3_NL/4)] * kvalues_iq3nl[idx]; + } + } + sumf += d * sumi; + } + *s = sumf; +} + void ggml_vec_dot_iq4_nl_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) { assert(nrc == 1); UNUSED(nrc); diff --git a/ggml/src/ggml-quants.c b/ggml/src/ggml-quants.c index c5471ca4cd9..b911c6917a1 100644 --- a/ggml/src/ggml-quants.c +++ b/ggml/src/ggml-quants.c @@ -26,14 +26,15 @@ #define UNUSED GGML_UNUSED static inline int best_index_int8(int n, const int8_t * val, float x) { - if (x <= val[0]) return 0; - if (x >= val[n-1]) return n-1; - int ml = 0, mu = n-1; - while (mu-ml > 1) { - int mav = (ml+mu)/2; - if (x < val[mav]) mu = mav; else ml = mav; - } - return x - val[mu-1] < val[mu] - x ? mu-1 : mu; + // val is sorted, so x picks level i once it reaches the midpoint of val[i-1] and val[i]. + // Scale by 2 to keep both sides exact, then binary search with no branches. n must be a power of 2. + const float x2 = x + x; + int idx = 0; + for (int step = n >> 1; step > 0; step >>= 1) { + const int j = idx + step; + idx = (x2 >= (float)(val[j-1] + val[j])) ? j : idx; + } + return idx; } // reference implementation for deterministic creation of model files @@ -5035,6 +5036,7 @@ static void quantize_row_iq4_nl_impl(const int super_block_size, const int block } if (amax < GROUP_MAX_EPS) { scales[ib] = 0; + memset(Lb, 0, block_size); continue; } float d = ntry > 0 ? -max/values[0] : max/values[0]; @@ -5113,6 +5115,7 @@ static void quantize_row_iq4_nl_impl(const int super_block_size, const int block } } } else { + GGML_ASSERT(nvalues <= 4 || qh); const int n = super_block_size/4; for (int j = 0; j < n; ++j) { q4[j] = (L[j] & 3) | ((L[j+n] & 3) << 2) | ((L[j+2*n] & 3) << 4) | ((L[j+3*n] & 3) << 6); From f47604d301988e996c5acc68a50f40d8bda9c922 Mon Sep 17 00:00:00 2001 From: Ed Addario Date: Tue, 1 Sep 2026 00:14:11 +0200 Subject: [PATCH 21/21] Refactor best_index_int8() --- ggml/src/ggml-cuda/cpy-utils.cuh | 13 ++++++------- ggml/src/ggml-metal/kernels/common.h | 13 ++++++------- ggml/src/ggml-sycl/cpy.hpp | 21 ++++++--------------- 3 files changed, 18 insertions(+), 29 deletions(-) diff --git a/ggml/src/ggml-cuda/cpy-utils.cuh b/ggml/src/ggml-cuda/cpy-utils.cuh index 7697c292dd6..ea4a09d8554 100644 --- a/ggml/src/ggml-cuda/cpy-utils.cuh +++ b/ggml/src/ggml-cuda/cpy-utils.cuh @@ -4,14 +4,13 @@ #include "convert.cuh" static __device__ __forceinline__ int best_index_int8(int n, const int8_t * val, float x) { - if (x <= val[0]) return 0; - if (x >= val[n-1]) return n-1; - int ml = 0, mu = n-1; - while (mu-ml > 1) { - int mav = (ml+mu)/2; - if (x < val[mav]) mu = mav; else ml = mav; + const float x2 = x + x; + int idx = 0; + for (int step = n >> 1; step > 0; step >>= 1) { + const int j = idx + step; + idx = (x2 >= (float)(val[j-1] + val[j])) ? j : idx; } - return x - val[mu-1] < val[mu] - x ? mu-1 : mu; + return idx; } static __device__ void quantize_f32_q4_0_block(const float * __restrict__ x, block_q4_0 * __restrict__ y) { diff --git a/ggml/src/ggml-metal/kernels/common.h b/ggml/src/ggml-metal/kernels/common.h index c4d67439448..0626715abcd 100644 --- a/ggml/src/ggml-metal/kernels/common.h +++ b/ggml/src/ggml-metal/kernels/common.h @@ -46,14 +46,13 @@ constexpr constant static float kvalues_mxfp4_f[16] = { }; static inline int best_index_int8(int n, constant float * val, float x) { - if (x <= val[0]) return 0; - if (x >= val[n-1]) return n-1; - int ml = 0, mu = n-1; - while (mu-ml > 1) { - int mav = (ml+mu)/2; - if (x < val[mav]) mu = mav; else ml = mav; + const float x2 = x + x; + int idx = 0; + for (int step = n >> 1; step > 0; step >>= 1) { + const int j = idx + step; + idx = (x2 >= val[j-1] + val[j]) ? j : idx; } - return x - val[mu-1] < val[mu] - x ? mu-1 : mu; + return idx; } static inline float e8m0_to_fp32(uint8_t x) { diff --git a/ggml/src/ggml-sycl/cpy.hpp b/ggml/src/ggml-sycl/cpy.hpp index 34bae1b2dd1..56d3d521ab8 100644 --- a/ggml/src/ggml-sycl/cpy.hpp +++ b/ggml/src/ggml-sycl/cpy.hpp @@ -7,22 +7,13 @@ typedef void (*cpy_kernel_t)(const char * cx, char * cdst); __dpct_inline__ int best_index_int8(int n, const int8_t * val, float x) { - if (x <= val[0]) { - return 0; - } - if (x >= val[n - 1]) { - return n - 1; - } - int ml = 0, mu = n - 1; - while (mu - ml > 1) { - int mav = (ml + mu) / 2; - if (x < val[mav]) { - mu = mav; - } else { - ml = mav; - } + const float x2 = x + x; + int idx = 0; + for (int step = n >> 1; step > 0; step >>= 1) { + const int j = idx + step; + idx = (x2 >= (float) (val[j - 1] + val[j])) ? j : idx; } - return x - val[mu - 1] < val[mu] - x ? mu - 1 : mu; + return idx; } inline void cpy_blck_f32_q8_0(const char * cxi, char * cdsti) {