diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h index 5f6774a630c..ce8a95b77f0 100644 --- a/ggml/include/ggml.h +++ b/ggml/include/ggml.h @@ -430,7 +430,9 @@ extern "C" { GGML_TYPE_NVFP4 = 40, // NVFP4 (4 blocks, E4M3 scale) GGML_TYPE_Q1_0 = 41, GGML_TYPE_Q2_0 = 42, - GGML_TYPE_COUNT = 43, + GGML_TYPE_IQ2_NL = 43, + GGML_TYPE_IQ3_NL = 44, + GGML_TYPE_COUNT = 45, }; // precision diff --git a/ggml/src/ggml-common.h b/ggml/src/ggml-common.h index 83f9118da84..b2724f928c0 100644 --- a/ggml/src/ggml-common.h +++ b/ggml/src/ggml-common.h @@ -444,6 +444,21 @@ typedef union { } iq1m_scale_t; // Non-linear quants +#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"); + #define QK4_NL 32 typedef struct { ggml_half d; @@ -1116,6 +1131,14 @@ GGML_TABLE_BEGIN(uint32_t, iq3s_grid, 512) 0x0f090307, 0x0f090501, 0x0f090b01, 0x0f0b0505, 0x0f0b0905, 0x0f0d0105, 0x0f0d0703, 0x0f0f0101, GGML_TABLE_END() +GGML_TABLE_BEGIN(int8_t, kvalues_iq2nl, 4) + -127, -49, 9, 88, +GGML_TABLE_END() + +GGML_TABLE_BEGIN(int8_t, kvalues_iq3nl, 8) + -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, diff --git a/ggml/src/ggml-cpu/arch-fallback.h b/ggml/src/ggml-cpu/arch-fallback.h index 152e0bac99b..8917666ddba 100644 --- a/ggml/src/ggml-cpu/arch-fallback.h +++ b/ggml/src/ggml-cpu/arch-fallback.h @@ -32,6 +32,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 @@ -119,6 +121,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 @@ -166,6 +170,8 @@ #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_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_0_4x8_generic ggml_quantize_mat_q8_0_4x8 @@ -207,6 +213,8 @@ // quants.c #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0 #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_4x1_generic ggml_quantize_mat_q8_0_4x1 #define ggml_quantize_mat_q8_0_4x4_generic ggml_quantize_mat_q8_0_4x4 @@ -259,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 +317,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 #define ggml_vec_dot_mxfp4_q8_0_generic ggml_vec_dot_mxfp4_q8_0 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-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c index b9c0fa3ddc0..042d3654691 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -376,6 +376,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, diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index b47ce5463c6..5a67cdadf60 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -685,6 +685,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: @@ -1137,6 +1139,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: @@ -1268,6 +1272,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: @@ -4537,6 +4543,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: @@ -4814,6 +4822,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: @@ -5039,6 +5049,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: @@ -5796,6 +5808,8 @@ 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: diff --git a/ggml/src/ggml-cpu/quants.c b/ggml/src/ggml-cpu/quants.c index 5e36459f8cb..69d64aa1d1d 100644 --- a/ggml/src/ggml-cpu/quants.c +++ b/ggml/src/ggml-cpu/quants.c @@ -1280,6 +1280,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); @@ -1328,6 +1391,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 93ea7eeffe5..7f94eee911b 100644 --- a/ggml/src/ggml-cpu/quants.h +++ b/ggml/src/ggml-cpu/quants.h @@ -34,6 +34,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); @@ -64,6 +66,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); @@ -98,6 +102,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); 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/ggml-metal-device.m b/ggml/src/ggml-metal/ggml-metal-device.m index 41ce90dc8a9..393712b6a26 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.m +++ b/ggml/src/ggml-metal/ggml-metal-device.m @@ -788,6 +788,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]; } @@ -1712,7 +1713,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: @@ -1776,7 +1780,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_F16) { 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-quants.c b/ggml/src/ggml-quants.c index 1ebc50a763f..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 @@ -2722,6 +2723,43 @@ void dequantize_row_iq1_m(const block_iq1_m * GGML_RESTRICT x, float * GGML_REST } } +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_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; @@ -4963,8 +5001,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, @@ -4974,7 +5013,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; @@ -4996,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]; @@ -5003,7 +5044,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]; @@ -5017,7 +5058,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]; @@ -5048,7 +5089,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; @@ -5062,14 +5103,29 @@ 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 { + 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); + } + 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; + } } } } @@ -5087,8 +5143,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); @@ -5107,8 +5163,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); } } @@ -5123,8 +5179,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); @@ -5137,6 +5193,81 @@ 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; + 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) { @@ -5645,6 +5776,14 @@ 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_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_IQ4_NL: { VALIDATE_ROW_DATA_D_F16_IMPL(block_iq4_nl, data, nb); diff --git a/ggml/src/ggml-quants.h b/ggml/src/ggml-quants.h index 75188f1af18..8c6ddb9243c 100644 --- a/ggml/src/ggml-quants.h +++ b/ggml/src/ggml-quants.h @@ -37,6 +37,8 @@ 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_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); @@ -71,6 +73,8 @@ 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_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); @@ -82,6 +86,8 @@ 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_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); 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) { diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index dc8a1744323..b762ae446f8 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -6004,6 +6004,11 @@ 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; } diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index e0b615c07ed..ec1a6950e45 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -860,6 +860,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, @@ -8002,6 +8018,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: diff --git a/gguf-py/gguf/constants.py b/gguf-py/gguf/constants.py index c99feb3c795..14b4a1cff8c 100644 --- a/gguf-py/gguf/constants.py +++ b/gguf-py/gguf/constants.py @@ -5511,6 +5511,8 @@ class GGMLQuantizationType(IntEnum): NVFP4 = 40 Q1_0 = 41 Q2_0 = 42 + IQ2_NL = 43 + IQ3_NL = 44 class ExpertGatingFuncType(IntEnum): @@ -5567,6 +5569,8 @@ class LlamaFileType(IntEnum): MOSTLY_NVFP4 = 39 # except 1d tensors MOSTLY_Q1_0 = 40 # except 1d tensors MOSTLY_Q2_0 = 41 # except 1d tensors + MOSTLY_IQ2_NL = 42 # except 1d tensors + MOSTLY_IQ3_NL = 43 # except 1d tensors GUESSED = 1024 # not specified in the model file @@ -5703,6 +5707,8 @@ class VisionProjectorType: GGMLQuantizationType.NVFP4: (64, 4 + 32), GGMLQuantizationType.Q1_0: (128, 2 + 16), GGMLQuantizationType.Q2_0: (64, 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..cdf278fdefd 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, -49, 9, 88) + + @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, -78, -45, -18, 3, 30, 62, 107) + + @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) diff --git a/include/llama.h b/include/llama.h index ef7a012c43a..3b132037f40 100644 --- a/include/llama.h +++ b/include/llama.h @@ -156,6 +156,8 @@ extern "C" { LLAMA_FTYPE_MOSTLY_NVFP4 = 39, // except 1d tensors LLAMA_FTYPE_MOSTLY_Q1_0 = 40, // except 1d tensors LLAMA_FTYPE_MOSTLY_Q2_0 = 41, // except 1d tensors + LLAMA_FTYPE_MOSTLY_IQ2_NL = 42, // except 1d tensors + LLAMA_FTYPE_MOSTLY_IQ3_NL = 43, // 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 1b1f852a010..b94b108ba13 100644 --- a/src/llama-model-loader.cpp +++ b/src/llama-model-loader.cpp @@ -66,6 +66,8 @@ const char * llama_ftype_name(llama_ftype ftype) { case LLAMA_FTYPE_MOSTLY_IQ3_XXS: name = LLAMA_FTYPE_PREFIX "IQ3_XXS - 3.0625 bpw"; break; case LLAMA_FTYPE_MOSTLY_IQ1_S: name = LLAMA_FTYPE_PREFIX "IQ1_S - 1.5625 bpw"; break; case LLAMA_FTYPE_MOSTLY_IQ1_M: name = LLAMA_FTYPE_PREFIX "IQ1_M - 1.75 bpw"; break; + case LLAMA_FTYPE_MOSTLY_IQ2_NL: name = LLAMA_FTYPE_PREFIX "IQ2_NL - 2.5 bpw"; break; + case LLAMA_FTYPE_MOSTLY_IQ3_NL: name = LLAMA_FTYPE_PREFIX "IQ3_NL - 3.5 bpw"; break; case LLAMA_FTYPE_MOSTLY_IQ4_NL: name = LLAMA_FTYPE_PREFIX "IQ4_NL - 4.5 bpw"; break; case LLAMA_FTYPE_MOSTLY_IQ4_XS: name = LLAMA_FTYPE_PREFIX "IQ4_XS - 4.25 bpw"; break; case LLAMA_FTYPE_MOSTLY_IQ3_S: name = LLAMA_FTYPE_PREFIX "IQ3_S - 3.4375 bpw"; break; @@ -766,6 +768,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; @@ -1059,12 +1063,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; diff --git a/src/llama-quant.cpp b/src/llama-quant.cpp index c414caa173f..bcb3c6a0b0a 100644 --- a/src/llama-quant.cpp +++ b/src/llama-quant.cpp @@ -385,10 +385,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_0: case GGML_TYPE_Q2_K: @@ -466,6 +466,12 @@ 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_IQ1_M) { @@ -499,6 +505,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; } @@ -509,8 +518,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) { @@ -526,8 +535,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)) { @@ -549,30 +558,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_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; @@ -590,9 +603,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; @@ -614,13 +627,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; } @@ -635,29 +654,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) { - 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()); @@ -806,9 +827,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; @@ -852,8 +875,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; @@ -861,6 +886,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/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp index 4a7a0623174..6e56881a780 100644 --- a/tests/test-backend-ops.cpp +++ b/tests/test-backend-ops.cpp @@ -2439,9 +2439,9 @@ struct test_set_rows : public test_case { } double max_nmse_err() override { - if (type_dst == GGML_TYPE_Q2_0 || 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_Q2_0 || 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 [-1,1], so it'll be // roughly (2.0 / 2^bits)^2, divided by the mean square value of the reference, @@ -2450,6 +2450,15 @@ struct test_set_rows : public test_case { if (type_src == GGML_TYPE_F16 && type_dst == GGML_TYPE_Q2_0) { err_estimate *= 4.0f; } + if (type_dst == 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_dst == 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_dst == GGML_TYPE_Q5_0 || type_dst == GGML_TYPE_Q5_1) { err_estimate /= 2.0f; } @@ -2973,13 +2982,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; @@ -8264,7 +8282,8 @@ static const ggml_type all_types[] = { // GGML_TYPE_TQ1_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[] = { @@ -8292,7 +8311,8 @@ static const ggml_type other_types[] = { // GGML_TYPE_TQ1_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, }; diff --git a/tests/test-quantize-fns.cpp b/tests/test-quantize-fns.cpp index 9510ac14ce0..79bd74c9a1c 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; @@ -161,8 +162,10 @@ static int test_vec_dot_q(bool verbose) { type == GGML_TYPE_Q2_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); @@ -179,8 +182,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 diff --git a/tools/quantize/quantize.cpp b/tools/quantize/quantize.cpp index 38950036cd8..4b5d3c6f831 100644 --- a/tools/quantize/quantize.cpp +++ b/tools/quantize/quantize.cpp @@ -57,6 +57,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", },