Skip to content

ggml: add q1_0_g128 quantization type #21766

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
oda-shuya wants to merge 1 commit into from
Closed

ggml: add q1_0_g128 quantization type #21766

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 3 additions & 1 deletion ggml/include/ggml.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Comment on lines +430 to +431

Copilot AI Apr 11, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The enum introduces a hole at value 40 (MXFP4 = 39 then Q1_0_g128 = 41), which makes GGML_TYPE_COUNT (42) larger than the number of actually-defined type IDs in this enum excerpt. Code that iterates for (t = 0; t < GGML_TYPE_COUNT; ++t) may now hit an undefined/unused type value (40) with zero-initialized traits, causing incorrect behavior. Recommended fix: assign GGML_TYPE_Q1_0_g128 = 40 and GGML_TYPE_COUNT = 41, unless 40 is intentionally reserved—if it is reserved, add an explicit placeholder enum value/comment and ensure the traits/dispatch logic handle that reserved ID safely.

Suggested change
GGML_TYPE_Q1_0_g128 = 41,
GGML_TYPE_COUNT = 42,
GGML_TYPE_Q1_0_g128 = 40,
GGML_TYPE_COUNT = 41,

Copilot uses AI. Check for mistakes.
};

// precision
Expand Down Expand Up @@ -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

Copilot AI Apr 11, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Like the ggml_type enum, this introduces a gap (25 then 27) without an obvious reserved value at 26. If ftype IDs are persisted (e.g., in file formats / metadata) or assumed contiguous anywhere, gaps can cause compatibility or iteration issues. Recommended fix: either use the next sequential value (= 26) or add/document an explicit reserved = 26 entry so the gap is intentional and discoverable.

Suggested change
GGML_FTYPE_MOSTLY_MXFP4 = 25, // except 1d tensors
GGML_FTYPE_MOSTLY_MXFP4 = 25, // except 1d tensors
GGML_FTYPE_RESERVED_26 = 26, // reserved to keep persisted ftype ids stable

Copilot uses AI. Check for mistakes.
GGML_FTYPE_MOSTLY_Q1_0_g128 = 27, // except 1d tensors
};

// available tensor operations:
Expand Down
7 changes: 7 additions & 0 deletions ggml/src/ggml-common.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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");
Comment on lines +239 to +244

Copilot AI Apr 11, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

New public identifiers mix uppercase and lowercase (QK1_0_g128, block_q1_0_g128, GGML_TYPE_Q1_0_g128). If the project convention is all-uppercase for macros/enums, consider renaming to an all-uppercase suffix (e.g., QK1_0_G128, GGML_TYPE_Q1_0_G128) while keeping the serialized/string name (\"q1_0_g128\") unchanged. This reduces inconsistency and makes grep/search across types more predictable.

Suggested change
#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");
#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");

Copilot uses AI. Check for mistakes.

//
// Ternary quantization
//
Expand Down
6 changes: 6 additions & 0 deletions ggml/src/ggml-cpu/ggml-cpu.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down
7 changes: 7 additions & 0 deletions ggml/src/ggml-cpu/ops.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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:
Expand Down Expand Up @@ -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:
Expand Down Expand Up @@ -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:
Expand Down Expand Up @@ -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:
Expand Down Expand Up @@ -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:
Expand Down Expand Up @@ -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:
Expand Down Expand Up @@ -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:
Expand Down
125 changes: 125 additions & 0 deletions ggml/src/ggml-cpu/quants.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
}
Expand Down Expand Up @@ -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);
Comment on lines +339 to +345

Copilot AI Apr 11, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The only optimized path here is __ARM_NEON; on non-NEON platforms (notably x86_64), the vec-dot falls back to the generic bit-extraction inner loop, which is likely to be a significant bottleneck relative to existing quant types with SIMD kernels. Consider adding an x86 SIMD implementation (SSE/AVX/AVX2/AVX-512 depending on existing project patterns) or gating/dispatching this type so it’s not selected for performance-critical CPU inference on platforms without a specialized kernel.

Copilot uses AI. Check for mistakes.
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);
Expand Down
3 changes: 3 additions & 0 deletions ggml/src/ggml-cpu/quants.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);

Expand All @@ -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);

Expand Down Expand Up @@ -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);

Expand Down
49 changes: 49 additions & 0 deletions ggml/src/ggml-quants.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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));
Comment on lines +225 to +237

Copilot AI Apr 11, 2026

Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The quantization scheme here stores only per-block sign bits plus a single shared magnitude d = max(abs(x)), meaning dequantization produces outputs in {+d, -d} (and cannot preserve per-element magnitudes). This is a very specific/atypical quantization behavior with significant error characteristics. Please add a short doc comment near the block definition and/or here explaining the intended use-case and the exact dequantized value set, so future kernel authors and users understand the tradeoffs and don’t assume it behaves like higher-bit quantizers.

Copilot uses AI. Check for mistakes.

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);
Expand Down Expand Up @@ -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;

Expand Down Expand Up @@ -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);
Expand Down Expand Up @@ -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);
Expand Down
3 changes: 3 additions & 0 deletions ggml/src/ggml-quants.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);

Expand All @@ -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);
Expand Down Expand Up @@ -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);

Expand Down
10 changes: 10 additions & 0 deletions ggml/src/ggml.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down Expand Up @@ -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;
Expand Down Expand Up @@ -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;
Expand Down
Loading