diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h index 77af0e7fb6aa..66522486ded7 100644 --- a/ggml/include/ggml.h +++ b/ggml/include/ggml.h @@ -427,7 +427,8 @@ extern "C" { // GGML_TYPE_IQ4_NL_4_8 = 37, // GGML_TYPE_IQ4_NL_8_8 = 38, GGML_TYPE_MXFP4 = 39, // MXFP4 (1 block) - GGML_TYPE_COUNT = 40, + GGML_TYPE_Q1_0_g128 = 41, + GGML_TYPE_COUNT = 42, }; // precision @@ -463,6 +464,7 @@ extern "C" { GGML_FTYPE_MOSTLY_IQ1_M = 23, // except 1d tensors GGML_FTYPE_MOSTLY_BF16 = 24, // except 1d tensors GGML_FTYPE_MOSTLY_MXFP4 = 25, // except 1d tensors + GGML_FTYPE_MOSTLY_Q1_0_g128 = 27, // except 1d tensors }; // available tensor operations: diff --git a/ggml/src/ggml-common.h b/ggml/src/ggml-common.h index 93ab7ea446e2..b3b7c0c447d3 100644 --- a/ggml/src/ggml-common.h +++ b/ggml/src/ggml-common.h @@ -236,6 +236,13 @@ typedef struct { } block_q8_1; static_assert(sizeof(block_q8_1) == 2*sizeof(ggml_half) + QK8_1, "wrong q8_1 block size/padding"); +#define QK1_0_g128 128 +typedef struct { + ggml_half d; // shared scale + uint8_t qs[QK1_0_g128/8]; // sign bits +} block_q1_0_g128; +static_assert(sizeof(block_q1_0_g128) == sizeof(ggml_half) + QK1_0_g128/8, "wrong q1_0_g128 block size/padding"); + // // Ternary quantization // diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c index 64eb01a4e18c..7668a2a54eb7 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -217,6 +217,12 @@ static const struct ggml_type_traits_cpu type_traits_cpu[GGML_TYPE_COUNT] = { .vec_dot_type = GGML_TYPE_F16, .nrows = 1, }, + [GGML_TYPE_Q1_0_g128] = { + .from_float = quantize_row_q1_0_g128, + .vec_dot = ggml_vec_dot_q1_0_g128_q8_0, + .vec_dot_type = GGML_TYPE_Q8_0, + .nrows = 1, + }, [GGML_TYPE_Q4_0] = { .from_float = quantize_row_q4_0, .vec_dot = ggml_vec_dot_q4_0_q8_0, diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index b7a70e06f1d0..3f5d396381dc 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -669,6 +669,7 @@ void ggml_compute_forward_add( case GGML_TYPE_Q5_0: case GGML_TYPE_Q5_1: case GGML_TYPE_Q8_0: + case GGML_TYPE_Q1_0_g128: case GGML_TYPE_MXFP4: case GGML_TYPE_Q2_K: case GGML_TYPE_Q3_K: @@ -1117,6 +1118,7 @@ void ggml_compute_forward_add1( case GGML_TYPE_Q5_0: case GGML_TYPE_Q5_1: case GGML_TYPE_Q8_0: + case GGML_TYPE_Q1_0_g128: case GGML_TYPE_Q8_1: case GGML_TYPE_MXFP4: case GGML_TYPE_Q2_K: @@ -1245,6 +1247,7 @@ void ggml_compute_forward_acc( case GGML_TYPE_Q5_0: case GGML_TYPE_Q5_1: case GGML_TYPE_Q8_0: + case GGML_TYPE_Q1_0_g128: case GGML_TYPE_Q8_1: case GGML_TYPE_MXFP4: case GGML_TYPE_Q2_K: @@ -4333,6 +4336,7 @@ void ggml_compute_forward_out_prod( case GGML_TYPE_Q5_0: case GGML_TYPE_Q5_1: case GGML_TYPE_Q8_0: + case GGML_TYPE_Q1_0_g128: case GGML_TYPE_MXFP4: case GGML_TYPE_Q2_K: case GGML_TYPE_Q3_K: @@ -4607,6 +4611,7 @@ void ggml_compute_forward_set( case GGML_TYPE_Q5_0: case GGML_TYPE_Q5_1: case GGML_TYPE_Q8_0: + case GGML_TYPE_Q1_0_g128: case GGML_TYPE_Q8_1: case GGML_TYPE_MXFP4: case GGML_TYPE_Q2_K: @@ -4829,6 +4834,7 @@ void ggml_compute_forward_get_rows( case GGML_TYPE_Q5_0: case GGML_TYPE_Q5_1: case GGML_TYPE_Q8_0: + case GGML_TYPE_Q1_0_g128: case GGML_TYPE_Q8_1: case GGML_TYPE_MXFP4: case GGML_TYPE_Q2_K: @@ -5553,6 +5559,7 @@ void ggml_compute_forward_clamp( case GGML_TYPE_Q5_0: case GGML_TYPE_Q5_1: case GGML_TYPE_Q8_0: + case GGML_TYPE_Q1_0_g128: case GGML_TYPE_Q8_1: case GGML_TYPE_MXFP4: case GGML_TYPE_Q2_K: diff --git a/ggml/src/ggml-cpu/quants.c b/ggml/src/ggml-cpu/quants.c index 365cb36d2d76..2728d96fd4e4 100644 --- a/ggml/src/ggml-cpu/quants.c +++ b/ggml/src/ggml-cpu/quants.c @@ -38,6 +38,10 @@ void quantize_row_q5_1(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, in quantize_row_q5_1_ref(x, y, k); } +void quantize_row_q1_0_g128(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { + quantize_row_q1_0_g128_ref(x, y, k); +} + void quantize_row_q8_0_generic(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { quantize_row_q8_0_ref(x, y, k); } @@ -332,6 +336,127 @@ void ggml_vec_dot_q8_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, c *s = sumf; } +void ggml_vec_dot_q1_0_g128_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) { + static_assert(QK1_0_g128 % QK8_0 == 0, "Q1_0_g128 must align with Q8_0 blocks"); + + const int blocks_per_q1 = QK1_0_g128 / QK8_0; + const int nb = n / QK1_0_g128; + + assert(n % QK1_0_g128 == 0); + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + + const block_q1_0_g128 * GGML_RESTRICT x = vx; + const block_q8_0 * GGML_RESTRICT y = vy; + + float sumf = 0.0f; + + for (int ib = 0; ib < nb; ++ib) { + float block_sum = 0.0f; + + for (int jb = 0; jb < blocks_per_q1; ++jb) { + const block_q8_0 * yb = &y[ib*blocks_per_q1 + jb]; + int sumi = 0; + + for (int j = 0; j < QK8_0; ++j) { + const int idx = jb*QK8_0 + j; + const int sign = ((x[ib].qs[idx / 8] >> (idx % 8)) & 1) ? 1 : -1; + sumi += sign * yb->qs[j]; + } + + block_sum += GGML_CPU_FP16_TO_FP32(yb->d) * (float) sumi; + } + + sumf += GGML_CPU_FP16_TO_FP32(x[ib].d) * block_sum; + } + + *s = sumf; +} + +#if defined(__ARM_NEON) +static const uint8_t k_q1_0_g128_bit_masks[8] = { 1, 2, 4, 8, 16, 32, 64, 128 }; + +static inline int8x8_t ggml_q1_0_g128_signs_neon(uint8_t packed_signs) { + const uint8x8_t sign_bits = vcgt_u8( + vand_u8(vdup_n_u8(packed_signs), vld1_u8(k_q1_0_g128_bit_masks)), + vdup_n_u8(0) + ); + return vreinterpret_s8_u8( + vbsl_u8( + sign_bits, + vreinterpret_u8_s8(vdup_n_s8(1)), + vreinterpret_u8_s8(vdup_n_s8(-1)) + ) + ); +} + +static inline int32_t ggml_vec_dot_q1_0_g128_q8_0_neon_16(int8x16_t signs, int8x16_t activations) { +#if defined(__ARM_FEATURE_DOTPROD) + return vaddvq_s32(ggml_vdotq_s32(vdupq_n_s32(0), signs, activations)); +#else + const int16x8_t prod_lo = vmull_s8(vget_low_s8(signs), vget_low_s8(activations)); + const int16x8_t prod_hi = vmull_s8(vget_high_s8(signs), vget_high_s8(activations)); + return vaddlvq_s16(vaddq_s16(prod_lo, prod_hi)); +#endif +} +#endif + +void ggml_vec_dot_q1_0_g128_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) { +#if defined(__ARM_NEON) + static_assert(QK1_0_g128 % QK8_0 == 0, "Q1_0_g128 must align with Q8_0 blocks"); + + const int blocks_per_q1 = QK1_0_g128 / QK8_0; + const int nb = n / QK1_0_g128; + + assert(n % QK1_0_g128 == 0); + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + + const block_q1_0_g128 * GGML_RESTRICT x = vx; + const block_q8_0 * GGML_RESTRICT y = vy; + + float sumf = 0.0f; + + for (int ib = 0; ib < nb; ++ib) { + float block_sum = 0.0f; + + for (int jb = 0; jb < blocks_per_q1; ++jb) { + const block_q8_0 * yb = &y[ib*blocks_per_q1 + jb]; + const uint8_t * packed_signs = x[ib].qs + jb*(QK8_0 / 8); + + const int8x16_t signs0 = vcombine_s8( + ggml_q1_0_g128_signs_neon(packed_signs[0]), + ggml_q1_0_g128_signs_neon(packed_signs[1]) + ); + const int8x16_t signs1 = vcombine_s8( + ggml_q1_0_g128_signs_neon(packed_signs[2]), + ggml_q1_0_g128_signs_neon(packed_signs[3]) + ); + const int8x16_t qy0 = vld1q_s8(yb->qs); + const int8x16_t qy1 = vld1q_s8(yb->qs + 16); + + const int sumi = + ggml_vec_dot_q1_0_g128_q8_0_neon_16(signs0, qy0) + + ggml_vec_dot_q1_0_g128_q8_0_neon_16(signs1, qy1); + + block_sum += GGML_CPU_FP16_TO_FP32(yb->d) * (float) sumi; + } + + sumf += GGML_CPU_FP16_TO_FP32(x[ib].d) * block_sum; + } + + *s = sumf; +#else + ggml_vec_dot_q1_0_g128_q8_0_generic(n, s, bs, vx, bx, vy, by, nrc); +#endif +} + void ggml_vec_dot_tq1_0_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); diff --git a/ggml/src/ggml-cpu/quants.h b/ggml/src/ggml-cpu/quants.h index d83eb1b144d4..a7ffaa693a3f 100644 --- a/ggml/src/ggml-cpu/quants.h +++ b/ggml/src/ggml-cpu/quants.h @@ -18,6 +18,7 @@ void quantize_row_q5_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, in void quantize_row_q5_1(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q8_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q8_1(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_q1_0_g128(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_mxfp4(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); @@ -40,6 +41,7 @@ void ggml_vec_dot_q4_1_q8_1(int n, float * GGML_RESTRICT s, size_t bs, const voi void ggml_vec_dot_q5_0_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_q5_1_q8_1(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_q8_0_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_q1_0_g128_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_mxfp4_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); @@ -71,6 +73,7 @@ void ggml_vec_dot_q4_1_q8_1_generic(int n, float * GGML_RESTRICT s, size_t bs, c void ggml_vec_dot_q5_0_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_q5_1_q8_1_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_q8_0_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_q1_0_g128_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_mxfp4_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); diff --git a/ggml/src/ggml-quants.c b/ggml/src/ggml-quants.c index de5cbd75e868..5edb56eafe49 100644 --- a/ggml/src/ggml-quants.c +++ b/ggml/src/ggml-quants.c @@ -221,6 +221,29 @@ void quantize_row_q8_0_ref(const float * GGML_RESTRICT x, block_q8_0 * GGML_REST } } +// reference implementation for deterministic creation of model files +void quantize_row_q1_0_g128_ref(const float * GGML_RESTRICT x, block_q1_0_g128 * GGML_RESTRICT y, int64_t k) { + assert(k % QK1_0_g128 == 0); + const int nb = k / QK1_0_g128; + + for (int i = 0; i < nb; ++i) { + float amax = 0.0f; + + for (int j = 0; j < QK1_0_g128; ++j) { + amax = MAX(amax, fabsf(x[i*QK1_0_g128 + j])); + } + + y[i].d = GGML_FP32_TO_FP16(amax); + memset(y[i].qs, 0, sizeof(y[i].qs)); + + for (int j = 0; j < QK1_0_g128; ++j) { + if (x[i*QK1_0_g128 + j] >= 0.0f) { + y[i].qs[j / 8] |= (uint8_t) (1u << (j % 8)); + } + } + } +} + // reference implementation for deterministic creation of model files void quantize_row_q8_1_ref(const float * GGML_RESTRICT x, block_q8_1 * GGML_RESTRICT y, int64_t k) { assert(QK8_1 == 32); @@ -414,6 +437,21 @@ void dequantize_row_q8_0(const block_q8_0 * GGML_RESTRICT x, float * GGML_RESTRI } } +void dequantize_row_q1_0_g128(const block_q1_0_g128 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + assert(k % QK1_0_g128 == 0); + + const int nb = k / QK1_0_g128; + + for (int i = 0; i < nb; ++i) { + const float d = GGML_FP16_TO_FP32(x[i].d); + + for (int j = 0; j < QK1_0_g128; ++j) { + const int bit = (x[i].qs[j / 8] >> (j % 8)) & 1; + y[i*QK1_0_g128 + j] = bit ? d : -d; + } + } +} + void dequantize_row_mxfp4(const block_mxfp4 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { static const int qk = QK_MXFP4; @@ -2092,6 +2130,13 @@ size_t quantize_q8_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, return nrow * row_size; } +size_t quantize_q1_0_g128(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { + GGML_UNUSED(quant_weights); + const size_t row_size = ggml_row_size(GGML_TYPE_Q1_0_g128, n_per_row); + quantize_row_q1_0_g128_ref(src, dst, (int64_t) nrow*n_per_row); + return nrow * row_size; +} + size_t quantize_mxfp4(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { GGML_UNUSED(quant_weights); quantize_row_mxfp4_ref(src, dst, (int64_t)nrow*n_per_row); @@ -5221,6 +5266,10 @@ bool ggml_validate_row_data(enum ggml_type type, const void * data, size_t nbyte { VALIDATE_ROW_DATA_D_F16_IMPL(block_q8_0, data, nb); } break; + case GGML_TYPE_Q1_0_g128: + { + VALIDATE_ROW_DATA_D_F16_IMPL(block_q1_0_g128, data, nb); + } break; case GGML_TYPE_MXFP4: { VALIDATE_ROW_DATA_E_E8M0_IMPL(block_mxfp4, data, nb); diff --git a/ggml/src/ggml-quants.h b/ggml/src/ggml-quants.h index 3b688f31c214..1b03b884702b 100644 --- a/ggml/src/ggml-quants.h +++ b/ggml/src/ggml-quants.h @@ -20,6 +20,7 @@ GGML_API void quantize_row_q5_0_ref(const float * GGML_RESTRICT x, block_q5_0 * GGML_API void quantize_row_q5_1_ref(const float * GGML_RESTRICT x, block_q5_1 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_q8_0_ref(const float * GGML_RESTRICT x, block_q8_0 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_q8_1_ref(const float * GGML_RESTRICT x, block_q8_1 * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_q1_0_g128_ref(const float * GGML_RESTRICT x, block_q1_0_g128 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_mxfp4_ref(const float * GGML_RESTRICT x, block_mxfp4 * GGML_RESTRICT y, int64_t k); @@ -45,6 +46,7 @@ GGML_API void dequantize_row_q4_1(const block_q4_1 * GGML_RESTRICT x, float * GG GGML_API void dequantize_row_q5_0(const block_q5_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_q5_1(const block_q5_1 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_q8_0(const block_q8_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_q1_0_g128(const block_q1_0_g128 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); //GGML_API void dequantize_row_q8_1(const block_q8_1 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_mxfp4(const block_mxfp4 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); @@ -93,6 +95,7 @@ GGML_API size_t quantize_q4_1(const float * GGML_RESTRICT src, void * GGML_RESTR GGML_API size_t quantize_q5_0(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_q5_1(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_q8_0(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_q1_0_g128(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_mxfp4(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.c b/ggml/src/ggml.c index ed819eaa4c5d..92d40e9df08f 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -651,6 +651,14 @@ static const struct ggml_type_traits type_traits[GGML_TYPE_COUNT] = { .to_float = (ggml_to_float_t) ggml_fp16_to_fp32_row, .from_float_ref = (ggml_from_float_t) ggml_fp32_to_fp16_row, }, + [GGML_TYPE_Q1_0_g128] = { + .type_name = "q1_0_g128", + .blck_size = QK1_0_g128, + .type_size = sizeof(block_q1_0_g128), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_q1_0_g128, + .from_float_ref = (ggml_from_float_t) quantize_row_q1_0_g128_ref, + }, [GGML_TYPE_Q4_0] = { .type_name = "q4_0", .blck_size = QK4_0, @@ -1364,6 +1372,7 @@ enum ggml_type ggml_ftype_to_ggml_type(enum ggml_ftype ftype) { case GGML_FTYPE_MOSTLY_Q5_0: wtype = GGML_TYPE_Q5_0; break; case GGML_FTYPE_MOSTLY_Q5_1: wtype = GGML_TYPE_Q5_1; break; case GGML_FTYPE_MOSTLY_Q8_0: wtype = GGML_TYPE_Q8_0; break; + case GGML_FTYPE_MOSTLY_Q1_0_g128: wtype = GGML_TYPE_Q1_0_g128; break; case GGML_FTYPE_MOSTLY_MXFP4: wtype = GGML_TYPE_MXFP4; break; case GGML_FTYPE_MOSTLY_Q2_K: wtype = GGML_TYPE_Q2_K; break; case GGML_FTYPE_MOSTLY_Q3_K: wtype = GGML_TYPE_Q3_K; break; @@ -7568,6 +7577,7 @@ size_t ggml_quantize_chunk( case GGML_TYPE_Q5_0: result = quantize_q5_0(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q5_1: result = quantize_q5_1(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q8_0: result = quantize_q8_0(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_Q1_0_g128: result = quantize_q1_0_g128(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_MXFP4: result = quantize_mxfp4(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q2_K: result = quantize_q2_K(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q3_K: result = quantize_q3_K(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;