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

MXFP4_R8 (#2196) 

* Adding MXFP4_R8 with AVX2 implementation

* Also offline repack

* Add AVX512 implementation for MXFP4_R8
This commit is contained in:
Kawrakow 2026-07-28 08:03:59 +03:00 committed by GitHub
parent 8a27bef8d4
commit f0f6ae4bb0
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
11 changed files with 412 additions and 20 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
examples/quantize/quantize.cpp
View File

@ -29,6 +29,7 @@ static const std::vector<struct quant_option> QUANT_OPTIONS = {
{ "Q5_1", LLAMA_FTYPE_MOSTLY_Q5_1, " 4.70G, +0.0349 ppl @ LLaMA-v1-7B", },
{ "Q6_0", LLAMA_FTYPE_MOSTLY_Q6_0, " 6.5 bpw quantization", },
{ "MXFP4", LLAMA_FTYPE_MOSTLY_MXFP4, " 4.25 bpw 4-bit float quantization",},
{ "MXFP4_R8", LLAMA_FTYPE_MOSTLY_MXFP4_R8, " MXFP4 repacked", },
{ "IQ2_XXS", LLAMA_FTYPE_MOSTLY_IQ2_XXS, " 2.06 bpw quantization", },
{ "IQ2_XXS_R4",LLAMA_FTYPE_MOSTLY_IQ2_XXS_R4,"IQ2_XXS repacked", },
{ "IQ2_XS", LLAMA_FTYPE_MOSTLY_IQ2_XS, " 2.31 bpw quantization", },

2
ggml/include/ggml.h
View File

@ -484,6 +484,7 @@ extern "C" {
GGML_TYPE_IQ5_K_R4 = 340,
GGML_TYPE_IQ4_KS_R4 = 344,
GGML_TYPE_IQ5_KS_R4 = 352,
GGML_TYPE_MXFP4_R8 = 353,
GGML_TYPE_Q8_K_R16 = 397,
GGML_TYPE_Q8_KV_R8 = 398,
GGML_TYPE_Q8_K_R8 = 399,
@ -582,6 +583,7 @@ extern "C" {
GGML_FTYPE_MOSTLY_IQ5_K_R4 = 333, // except 1d tensors
GGML_FTYPE_MOSTLY_IQ4_KS_R4 = 337, // except 1d tensors
GGML_FTYPE_MOSTLY_IQ5_KS_R4 = 341, // except 1d tensors
GGML_FTYPE_MOSTLY_MXFP4_R8 = 342, // except 1d tensors
GGML_FTYPE_MOSTLY_Q8_K_R16 = 397, // except 1d tensors
GGML_FTYPE_MOSTLY_Q8_KV_R8 = 398, // except 1d tensors
GGML_FTYPE_MOSTLY_Q8_K_R8 = 399, // except 1d tensors

6
ggml/src/ggml-common.h
View File

@ -186,6 +186,12 @@ typedef struct {
} block_mxfp4;
static_assert(sizeof(block_mxfp4) == sizeof(uint8_t) + QK_MXFP4/2, "wrong mxfp4 block size/padding");
typedef struct {
uint8_t e[8]; // E8M0
uint8_t qs[4*QK_MXFP4];
} block_mxfp4_r8;
static_assert(sizeof(block_mxfp4_r8) == 8*sizeof(block_mxfp4), "wrong mxfp4_r8 block size/padding");
#define QK5_0 32
typedef struct {
ggml_half d; // delta

1
ggml/src/ggml-quants.c
View File

@ -15504,6 +15504,7 @@ bool ggml_validate_row_data(enum ggml_type type, const void * data, size_t nbyte
VALIDATE_ROW_DATA_D_F16_IMPL(block_iq4_nl, data, nb);
} break;
case GGML_TYPE_MXFP4: break;
case GGML_TYPE_MXFP4_R8: break;
case GGML_TYPE_Q6_0: break;
case GGML_TYPE_IQ2_K: break;
case GGML_TYPE_IQ2_KS: break;

26
ggml/src/ggml.c
View File

@ -1340,6 +1340,23 @@ static const ggml_type_traits_t type_traits[GGML_TYPE_COUNT] = {
.vec_dot_type = GGML_TYPE_Q8_2_X4,
#else
.vec_dot_type = GGML_TYPE_Q8_0_X4,
#endif
.nrows = 1,
.row_meta_size = 0,
},
[GGML_TYPE_MXFP4_R8] = {
.type_name = "mxfp4_r8",
.blck_size = QK_MXFP4,
.type_size = sizeof(block_mxfp4),
.is_quantized = true,
.to_float = (ggml_to_float_t) dequantize_row_mxfp4_r8,
.from_float = quantize_row_mxfp4_r8,
.from_float_ref = (ggml_from_float_t)quantize_row_mxfp4_r8_ref,
.vec_dot = vec_dot_mxfp4_r8_q8_2_x4,
#if defined __AVX2__
.vec_dot_type = GGML_TYPE_Q8_2_X4,
#else
.vec_dot_type = GGML_TYPE_Q8_0_X4,
#endif
.nrows = 1,
.row_meta_size = 0,
@ -4940,6 +4957,7 @@ enum ggml_type ggml_ftype_to_ggml_type(enum ggml_ftype ftype) {
case GGML_FTYPE_MOSTLY_Q8_0_R8: wtype = GGML_TYPE_Q8_0_R8; break;
case GGML_FTYPE_MOSTLY_IQ4_XS: wtype = GGML_TYPE_IQ4_XS; break;
case GGML_FTYPE_MOSTLY_MXFP4: wtype = GGML_TYPE_MXFP4; break;
case GGML_FTYPE_MOSTLY_MXFP4_R8: wtype = GGML_TYPE_MXFP4_R8; break;
case GGML_FTYPE_MOSTLY_IQ4_KS: wtype = GGML_TYPE_IQ4_KS; break;
case GGML_FTYPE_MOSTLY_IQ4_KS_R4: wtype = GGML_TYPE_IQ4_KS_R4;break;
case GGML_FTYPE_MOSTLY_IQ5_KS_R4: wtype = GGML_TYPE_IQ5_KS_R4;break;
@ -13434,6 +13452,7 @@ static void ggml_compute_forward_add(
case GGML_TYPE_I2_S:
case GGML_TYPE_Q8_0_R8:
case GGML_TYPE_MXFP4:
case GGML_TYPE_MXFP4_R8:
case GGML_TYPE_IQ4_XS:
case GGML_TYPE_IQ4_KS:
case GGML_TYPE_IQ4_KS_R4:
@ -13988,6 +14007,7 @@ static void ggml_compute_forward_add1(
case GGML_TYPE_I2_S:
case GGML_TYPE_Q8_0_R8:
case GGML_TYPE_MXFP4:
case GGML_TYPE_MXFP4_R8:
case GGML_TYPE_IQ4_XS:
case GGML_TYPE_IQ4_KS:
case GGML_TYPE_IQ4_KS_R4:
@ -14168,6 +14188,7 @@ static void ggml_compute_forward_acc(
case GGML_TYPE_I2_S:
case GGML_TYPE_Q8_0_R8:
case GGML_TYPE_MXFP4:
case GGML_TYPE_MXFP4_R8:
case GGML_TYPE_IQ4_XS:
case GGML_TYPE_IQ4_KS:
case GGML_TYPE_IQ4_KS_R4:
@ -18846,6 +18867,7 @@ static void ggml_compute_forward_out_prod(
case GGML_TYPE_I2_S:
case GGML_TYPE_Q8_0_R8:
case GGML_TYPE_MXFP4:
case GGML_TYPE_MXFP4_R8:
case GGML_TYPE_IQ4_XS:
case GGML_TYPE_IQ4_KS:
case GGML_TYPE_IQ4_KS_R4:
@ -19270,6 +19292,7 @@ static void ggml_compute_forward_set(
case GGML_TYPE_I2_S:
case GGML_TYPE_Q8_0_R8:
case GGML_TYPE_MXFP4:
case GGML_TYPE_MXFP4_R8:
case GGML_TYPE_IQ4_XS:
case GGML_TYPE_IQ4_KS:
case GGML_TYPE_IQ4_KS_R4:
@ -19600,6 +19623,7 @@ static void ggml_compute_forward_get_rows(
case GGML_TYPE_I2_S:
case GGML_TYPE_Q8_0_R8:
case GGML_TYPE_MXFP4:
case GGML_TYPE_MXFP4_R8:
case GGML_TYPE_IQ4_XS:
case GGML_TYPE_IQ4_KS:
case GGML_TYPE_IQ4_KS_R4:
@ -20367,6 +20391,7 @@ static void ggml_compute_forward_clamp(
case GGML_TYPE_I2_S:
case GGML_TYPE_Q8_0_R8:
case GGML_TYPE_MXFP4:
case GGML_TYPE_MXFP4_R8:
case GGML_TYPE_IQ4_XS:
case GGML_TYPE_IQ4_KS:
case GGML_TYPE_IQ4_KS_R4:
@ -30368,6 +30393,7 @@ size_t ggml_quantize_chunk(
case GGML_TYPE_Q6_0_R4: result = quantize_q6_0_r4(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix, user_data); break;
case GGML_TYPE_Q8_0_R8: result = quantize_q8_0_r8(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix, user_data); break;
case GGML_TYPE_MXFP4: result = quantize_mxfp4 (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix, user_data); break;
case GGML_TYPE_MXFP4_R8:result = quantize_mxfp4_r8(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix, user_data); break;
case GGML_TYPE_IQ4_XS: result = quantize_iq4_xs (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix, user_data); break;
case GGML_TYPE_IQ4_KS: result = quantize_iq4_ks (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix, user_data); break;
case GGML_TYPE_IQ4_KS_R4:result = quantize_iq4_ks_r4(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix, user_data); break;

250
ggml/src/iqk/iqk_gemm_legacy_quants.cpp
View File

@ -1193,6 +1193,250 @@ static void mul_mat_q4_0_r8_q8_2(int n, const void * vx, size_t bx, const DataIn
}
#endif
inline void prepare_mxfp4_quants_avx2(const uint8_t * qs, __m256i * v, const __m256i& m4, const __m256i & table) {
auto bits1 = _mm256_loadu_si256((const __m256i *)qs+0);
auto bits2 = _mm256_loadu_si256((const __m256i *)qs+1);
auto bits3 = _mm256_loadu_si256((const __m256i *)qs+2);
auto bits4 = _mm256_loadu_si256((const __m256i *)qs+3);
v[0] = _mm256_shuffle_epi8(table, _mm256_and_si256(bits1, m4));
v[1] = _mm256_shuffle_epi8(table, _mm256_and_si256(bits2, m4));
v[2] = _mm256_shuffle_epi8(table, _mm256_and_si256(bits3, m4));
v[3] = _mm256_shuffle_epi8(table, _mm256_and_si256(bits4, m4));
v[4] = _mm256_shuffle_epi8(table, _mm256_and_si256(_mm256_srli_epi16(bits1, 4), m4));
v[5] = _mm256_shuffle_epi8(table, _mm256_and_si256(_mm256_srli_epi16(bits2, 4), m4));
v[6] = _mm256_shuffle_epi8(table, _mm256_and_si256(_mm256_srli_epi16(bits3, 4), m4));
v[7] = _mm256_shuffle_epi8(table, _mm256_and_si256(_mm256_srli_epi16(bits4, 4), m4));
}
inline __m256i accum_mxfp4_quants(const __m256i * v, const int8_t * qs) {
auto y4l = _mm_loadu_si128((const __m128i*)qs+0);
auto y4h = _mm_loadu_si128((const __m128i*)qs+1);
auto yl = MM256_SET1_M128I(y4l);
auto yh = MM256_SET1_M128I(y4h);
#ifdef HAVE_VNNI256
auto sumi = _mm256_setzero_si256();
sumi = ggml_mm256_dpbusd_epi32(sumi, v[0], _mm256_shuffle_epi32(yl, 0x00));
sumi = ggml_mm256_dpbusd_epi32(sumi, v[1], _mm256_shuffle_epi32(yl, 0x55));
sumi = ggml_mm256_dpbusd_epi32(sumi, v[2], _mm256_shuffle_epi32(yl, 0xaa));
sumi = ggml_mm256_dpbusd_epi32(sumi, v[3], _mm256_shuffle_epi32(yl, 0xff));
sumi = ggml_mm256_dpbusd_epi32(sumi, v[4], _mm256_shuffle_epi32(yh, 0x00));
sumi = ggml_mm256_dpbusd_epi32(sumi, v[5], _mm256_shuffle_epi32(yh, 0x55));
sumi = ggml_mm256_dpbusd_epi32(sumi, v[6], _mm256_shuffle_epi32(yh, 0xaa));
sumi = ggml_mm256_dpbusd_epi32(sumi, v[7], _mm256_shuffle_epi32(yh, 0xff));
#else
auto sumi1 = _mm256_add_epi16(_mm256_maddubs_epi16(v[0], _mm256_shuffle_epi32(yl, 0x00)),
_mm256_maddubs_epi16(v[1], _mm256_shuffle_epi32(yl, 0x55)));
auto sumi2 = _mm256_add_epi16(_mm256_maddubs_epi16(v[2], _mm256_shuffle_epi32(yl, 0xaa)),
_mm256_maddubs_epi16(v[3], _mm256_shuffle_epi32(yl, 0xff)));
auto sumi3 = _mm256_add_epi16(_mm256_maddubs_epi16(v[4], _mm256_shuffle_epi32(yh, 0x00)),
_mm256_maddubs_epi16(v[5], _mm256_shuffle_epi32(yh, 0x55)));
auto sumi4 = _mm256_add_epi16(_mm256_maddubs_epi16(v[6], _mm256_shuffle_epi32(yh, 0xaa)),
_mm256_maddubs_epi16(v[7], _mm256_shuffle_epi32(yh, 0xff)));
auto m1 = _mm256_set1_epi16(1);
auto sumi12 = _mm256_add_epi32(_mm256_madd_epi16(m1, sumi1), _mm256_madd_epi16(m1, sumi2));
auto sumi34 = _mm256_add_epi32(_mm256_madd_epi16(m1, sumi3), _mm256_madd_epi16(m1, sumi4));
auto sumi = _mm256_add_epi32(sumi12, sumi34);
#endif
return sumi;
}
inline __m256 convert_mxfp4_scales(const uint8_t * e) {
auto aux = _mm256_cvtepu8_epi32(_mm_loadl_epi64((const __m128i *)e));
auto mask = _mm256_cmpgt_epi32(aux, _mm256_set1_epi32(1));
auto d1 = _mm256_slli_epi32(_mm256_sub_epi32(aux, _mm256_set1_epi32(1)), 23);
auto d2 = _mm256_slli_epi32(_mm256_add_epi32(aux, _mm256_set1_epi32(1)), 21);
return _mm256_castsi256_ps(_mm256_or_si256(_mm256_and_si256(mask, d1), _mm256_andnot_si256(mask, d2)));
}
template <int nrc_y>
static void mul_mat_mxfp4_r8_q8_2_avx2(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) {
GGML_ASSERT(nrc_x%8 == 0);
Q8<nrc_y, block_q8_2_x4> q8(info);
auto m4 = _mm256_set1_epi8(0xf);
int nb = n / QK_MXFP4;
auto table128 = _mm_loadu_si128((const __m128i *)kvalues_mxfp4);
auto table = MM256_SET1_M128I(table128);
table = _mm256_add_epi8(table, _mm256_set1_epi8(12));
__m256i v[8];
if constexpr (nrc_y == 1) {
union { __m256 vec; float val[8]; } helper;
for (int ix = 0; ix < nrc_x; ix += 8) {
auto * iq4 = (const block_mxfp4_r8 *)((const char *)vx + ix*bx);
auto acc1 = _mm256_setzero_ps();
auto acc2 = _mm256_setzero_ps();
for (int ib4 = 0; ib4 < nb/4; ++ib4) {
helper.vec = convert_scales((const uint16_t *)q8.y[0][ib4].d);
for (int k = 0; k < 4; ++k) {
auto scales = convert_mxfp4_scales(iq4[4*ib4+k].e);
prepare_mxfp4_quants_avx2(iq4[4*ib4+k].qs, v, m4, table);
auto sumi = accum_mxfp4_quants(v, q8.y[0][ib4].qs+32*k);
auto d4d8 = _mm256_mul_ps(scales, _mm256_set1_ps(helper.val[k]));
acc1 = _mm256_fmadd_ps(d4d8, _mm256_cvtepi32_ps(sumi), acc1);
acc2 = _mm256_fmadd_ps(scales, _mm256_set1_ps(helper.val[k+4]), acc2);
}
}
for (int ib = 4*(nb/4); ib < nb; ++ib) {
auto qy = (const block_q8_2 *)q8.y[0];
auto scales = convert_mxfp4_scales(iq4[ib].e);
prepare_mxfp4_quants_avx2(iq4[ib].qs, v, m4, table);
auto sumi = accum_mxfp4_quants(v, qy[ib].qs);
auto [d8, m8] = ScaleHelperQ8_2::prepare1(qy + ib);
auto d4d8 = _mm256_mul_ps(scales, _mm256_set1_ps(d8));
acc1 = _mm256_fmadd_ps(d4d8, _mm256_cvtepi32_ps(sumi), acc1);
acc2 = _mm256_fmadd_ps(scales, _mm256_set1_ps(m8), acc2);
}
acc1 = _mm256_fmadd_ps(acc2, _mm256_set1_ps(-12.f), acc1);
info.store(ix, 0, acc1);
}
}
else {
__m256 acc[nrc_y] = {};
float d8[8*nrc_y];
for (int ix = 0; ix < nrc_x; ix += 8) {
auto * iq4 = (const block_mxfp4_r8 *)((const char *)vx + ix*bx);
for (int ib4 = 0; ib4 < nb/4; ++ib4) {
__m256 d4[4];
{
for (int k = 0; k < 4; ++k) {
d4[k] = convert_mxfp4_scales(iq4[4*ib4+k].e);
}
for (int iy = 0; iy < nrc_y; ++iy) {
auto scales = convert_scales((const uint16_t *)q8.y[iy][ib4].d);
_mm256_storeu_ps(d8 + 8*iy, scales);
auto m4 = _mm256_extractf128_ps(scales, 1);
auto m8 = _mm256_set_m128(m4, m4);
auto sumf = _mm256_mul_ps(d4[0], _mm256_shuffle_ps(m8, m8, 0x00));
sumf = _mm256_fmadd_ps(d4[1], _mm256_shuffle_ps(m8, m8, 0x55), sumf);
sumf = _mm256_fmadd_ps(d4[2], _mm256_shuffle_ps(m8, m8, 0xaa), sumf);
sumf = _mm256_fmadd_ps(d4[3], _mm256_shuffle_ps(m8, m8, 0xff), sumf);
acc[iy] = _mm256_fmadd_ps(sumf, _mm256_set1_ps(-12.f), acc[iy]);
}
}
for (int k = 0; k < 4; ++k) {
//auto scales = _mm256_cvtph_ps(_mm_loadu_si128((const __m128i *)iq4[4*ib4+k].d));
prepare_mxfp4_quants_avx2(iq4[4*ib4+k].qs, v, m4, table);
for (int iy = 0; iy < nrc_y; ++iy) {
auto sumi = accum_mxfp4_quants(v, q8.y[iy][ib4].qs+32*k);
auto d4d8 = _mm256_mul_ps(d4[k], _mm256_set1_ps(d8[8*iy+k]));
acc[iy] = _mm256_fmadd_ps(d4d8, _mm256_cvtepi32_ps(sumi), acc[iy]);
}
}
}
for (int ib = 4*(nb/4); ib < nb; ++ib) {
auto scales = convert_mxfp4_scales(iq4[ib].e);
auto scales_m = _mm256_mul_ps(scales, _mm256_set1_ps(-12.f));
prepare_mxfp4_quants_avx2(iq4[ib].qs, v, m4, table);
for (int iy = 0; iy < nrc_y; ++iy) {
auto qy = (const block_q8_2 *)q8.y[iy];
auto sumi = accum_mxfp4_quants(v, qy[ib].qs);
auto [d8, m8] = ScaleHelperQ8_2::prepare1(qy + ib);
auto d4d8 = _mm256_mul_ps(scales, _mm256_set1_ps(d8));
acc[iy] = _mm256_fmadd_ps(d4d8, _mm256_cvtepi32_ps(sumi), acc[iy]);
acc[iy] = _mm256_fmadd_ps(scales_m, _mm256_set1_ps(m8), acc[iy]);
}
}
for (int iy = 0; iy < nrc_y; ++iy) {
info.store(ix, iy, acc[iy]);
acc[iy] = _mm256_setzero_ps();
}
}
}
}
#ifdef HAVE_FANCY_SIMD
template <int nrc_y>
static void mul_mat_mxfp4_r8_q8_2(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) {
if constexpr (nrc_y == 1) {
mul_mat_mxfp4_r8_q8_2_avx2<1>(n, vx, bx, info, nrc_x);
return;
}
GGML_ASSERT(nrc_x%16 == 0);
Q8<nrc_y, block_q8_2_x4> q8(info);
auto m4 = _mm512_set1_epi8(0xf);
int nb = n / QK4_NL;
//auto table = _mm512_broadcast_i32x4(_mm_loadu_si128((const __m128i *)kvalues_mxfp4));
auto table128 = _mm_loadu_si128((const __m128i *)kvalues_mxfp4);
auto table256 = MM256_SET1_M128I(table128);
auto table = _mm512_inserti32x8(_mm512_castsi256_si512(table256), table256, 1);
table = _mm512_add_epi8(table, _mm512_set1_epi8(12));
__m512 acc[2*nrc_y] = {};
__m512i qx[8];
auto prepare = [&qx, &m4, &table] (const block_mxfp4_r8& iq4l, const block_mxfp4_r8& iq4h) {
auto scales1 = convert_mxfp4_scales(iq4l.e);
auto scales2 = convert_mxfp4_scales(iq4h.e);
auto scales = _mm512_insertf32x8(_mm512_castps256_ps512(scales1), scales2, 1);
for (int j = 0; j < 4; ++j) {
auto bits = _mm512_inserti32x8(_mm512_castsi256_si512(_mm256_loadu_si256((const __m256i *)iq4l.qs+j)),
_mm256_loadu_si256((const __m256i *)iq4h.qs+j), 1);
qx[j+0] = _mm512_and_si512(bits, m4);
qx[j+4] = _mm512_and_si512(_mm512_srli_epi16(bits, 4), m4);
}
for (int j = 0; j < 8; ++j) qx[j] = _mm512_shuffle_epi8(table, qx[j]);
return scales;
};
auto dot = [&qx] (const int8_t * qy) {
auto y4l = _mm_loadu_si128((const __m128i*)qy+0);
auto y4h = _mm_loadu_si128((const __m128i*)qy+1);
//auto yl = _mm512_broadcast_i32x4(y4l);
//auto yh = _mm512_broadcast_i32x4(y4h);
auto y8l = MM256_SET1_M128I(y4l);
auto y8h = MM256_SET1_M128I(y4h);
auto yl = _mm512_inserti32x8(_mm512_castsi256_si512(y8l), y8l, 1);
auto yh = _mm512_inserti32x8(_mm512_castsi256_si512(y8h), y8h, 1);
auto sumi = _mm512_setzero_si512();
sumi = _mm512_dpbusd_epi32(sumi, qx[0], _mm512_shuffle_epi32(yl, _MM_PERM_ENUM(0x00)));
sumi = _mm512_dpbusd_epi32(sumi, qx[1], _mm512_shuffle_epi32(yl, _MM_PERM_ENUM(0x55)));
sumi = _mm512_dpbusd_epi32(sumi, qx[2], _mm512_shuffle_epi32(yl, _MM_PERM_ENUM(0xaa)));
sumi = _mm512_dpbusd_epi32(sumi, qx[3], _mm512_shuffle_epi32(yl, _MM_PERM_ENUM(0xff)));
sumi = _mm512_dpbusd_epi32(sumi, qx[4], _mm512_shuffle_epi32(yh, _MM_PERM_ENUM(0x00)));
sumi = _mm512_dpbusd_epi32(sumi, qx[5], _mm512_shuffle_epi32(yh, _MM_PERM_ENUM(0x55)));
sumi = _mm512_dpbusd_epi32(sumi, qx[6], _mm512_shuffle_epi32(yh, _MM_PERM_ENUM(0xaa)));
sumi = _mm512_dpbusd_epi32(sumi, qx[7], _mm512_shuffle_epi32(yh, _MM_PERM_ENUM(0xff)));
return sumi;
};
float d8[8*nrc_y];
for (int ix = 0; ix < nrc_x; ix += 16) {
auto iq4l = (const block_mxfp4_r8 *)((const char *)vx + (ix+0)*bx);
auto iq4h = (const block_mxfp4_r8 *)((const char *)vx + (ix+8)*bx);
for (int ib4 = 0; ib4 < nb/4; ++ib4) {
for (int iy = 0; iy < nrc_y; ++iy) {
_mm256_storeu_ps(d8+8*iy, convert_scales((const uint16_t *)q8.y[iy][ib4].d));
}
for (int k = 0; k < 4; ++k) {
auto scales = prepare(iq4l[4*ib4+k], iq4h[4*ib4+k]);
for (int iy = 0; iy < nrc_y; ++iy) {
auto sumi = dot(q8.y[iy][ib4].qs+32*k);
auto dy = _mm512_set1_ps(d8[8*iy+k]);
acc[2*iy+0] = _mm512_fmadd_ps(_mm512_mul_ps(scales, dy), _mm512_cvtepi32_ps(sumi), acc[2*iy+0]);
acc[2*iy+1] = _mm512_fmadd_ps(scales, _mm512_set1_ps(d8[8*iy+k+4]), acc[2*iy+1]);
}
}
}
for (int ib = 4*(nb/4); ib < nb; ++ib) {
auto scales = prepare(iq4l[ib], iq4h[ib]);
for (int iy = 0; iy < nrc_y; ++iy) {
auto qy = (const block_q8_1 *)q8.y[iy];
auto sumi = dot(qy[ib].qs);
auto [d8, m8] = ScaleHelperQ8_2::prepare1(qy + ib);
auto dy = _mm512_set1_ps(d8);
acc[2*iy+0] = _mm512_fmadd_ps(_mm512_mul_ps(scales, dy), _mm512_cvtepi32_ps(sumi), acc[2*iy+0]);
acc[2*iy+1] = _mm512_fmadd_ps(scales, _mm512_set1_ps(m8), acc[2*iy+1]);
}
}
for (int iy = 0; iy < nrc_y; ++iy) {
auto sum = _mm512_fmadd_ps(_mm512_set1_ps(-12.f), acc[2*iy+1], acc[2*iy+0]);
acc[2*iy+0] = acc[2*iy+1] = _mm512_setzero_ps();
info.store(ix, iy, sum);
}
}
}
#else
template <int nrc_y>
static void mul_mat_mxfp4_r8_q8_2(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) {
mul_mat_mxfp4_r8_q8_2_avx2<nrc_y>(n, vx, bx, info, nrc_x);
}
#endif
template <int nrc_y>
static void mul_mat_q5_0_r4_q8_2_avx2(int n, const void * vx, size_t bx, const DataInfo& info, int nrc_x) {
GGML_ASSERT(nrc_x%4 == 0);
@ -2128,6 +2372,12 @@ bool iqk_set_kernels_legacy_quants(int ne00, int typeA, int typeB, std::array<mu
func16 = mul_mat_q4_0_r8_q8_2<16>;
#endif
break;
case GGML_TYPE_MXFP4_R8:
IQK_SET_MUL_MAT_FUNCTIONS(mul_mat_mxfp4_r8_q8_2, kernels)
//#ifdef HAVE_FANCY_SIMD
// func16 = mul_mat_mxfp4_r8_q8_2<16>;
//#endif
break;
case GGML_TYPE_Q5_0_R4:
IQK_SET_MUL_MAT_FUNCTIONS(mul_mat_q5_0_r4_q8_2, kernels)
break;

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

@ -357,6 +357,7 @@ struct MulMat {
case GGML_TYPE_Q8_0_R8:
case GGML_TYPE_Q8_1:
case GGML_TYPE_Q8_K_R16:
case GGML_TYPE_MXFP4_R8:
case GGML_TYPE_BF16_R16: return 16;
default: return 1;
}
@ -390,6 +391,7 @@ struct MulMat {
case GGML_TYPE_Q8_KV:
case GGML_TYPE_Q8_KV_R8:
case GGML_TYPE_Q8_1:
case GGML_TYPE_MXFP4_R8:
case GGML_TYPE_Q8_K_R8: return 8;
case GGML_TYPE_Q8_K_R16:
case GGML_TYPE_BF16_R16: return 16;
@ -933,6 +935,7 @@ bool MulMat::prepare(int typeA, int typeB, int ne00, MulMat& mm, int Ny) {
case GGML_TYPE_Q8_0_R8:
case GGML_TYPE_IQ4_NL_R4:
case GGML_TYPE_MXFP4:
case GGML_TYPE_MXFP4_R8:
return iqk_set_kernels_legacy_quants(ne00, typeA, typeB, mm.funcs, mm.func16);
case GGML_TYPE_IQ1_S:
case GGML_TYPE_IQ1_M:

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

@ -4175,29 +4175,9 @@ static void quantize_row_mxfp4_impl(int n_per_row, const float * x, char * cy,
block_mxfp4 * y = (block_mxfp4 *)cy;
//int last_ibl = -1;
//float sigma2 = 0;
//const uint8_t e = (uint8_t) (floorf(log2f(amax)) - 2 + 127);
// -> log2f(amax) ~ e - 125 -> amax = 2^(e - 125)
//const float d = GGML_E8M0_TO_FP32_HALF(e);
for (int ib = 0; ib < n_per_row/QK_MXFP4; ++ib) {
memset(&y[ib], 0, sizeof(block_mxfp4));
const float * xb = x + ib*QK_MXFP4;
//if (int ibl = ib/(QK_K/QK_MXFP4); ibl != last_ibl) {
// int n = std::min(QK_K, n_per_row - ib*QK_MXFP4);
// float sumx2 = 0;
// for (int j = 0; j < n; ++j) sumx2 += xb[j]*xb[j];
// sigma2 = 2.0f*sumx2/n;
// last_ibl = ibl;
//}
//if (quant_weights) {
// const float * qw = quant_weights + ib*QK_MXFP4;
// for (int j = 0; j < QK_MXFP4; ++j) weight[j] = qw[j] * sqrtf(sigma2 + xb[j]*xb[j]);
//} else {
// for (int j = 0; j < QK_MXFP4; ++j) weight[j] = xb[j]*xb[j];
//}
float amax = 0;
for (int j = 0; j < QK_MXFP4; ++j) {
float ax = fabsf(xb[j]);
@ -4294,6 +4274,97 @@ void vec_dot_mxfp4_q8_0_x4(int n, float * s, size_t bs, const void * vx, size_t
//*s = sumf;
}
void quantize_row_mxfp4_r8_ref(const float * x, block_mxfp4_r8 * y, int64_t k) {
quantize_mxfp4_r8(x, (void *)y, 8, k/8, nullptr, nullptr);
}
void quantize_row_mxfp4_r8(const float * x, void * y, int64_t k) {
quantize_mxfp4_r8(x, (void *)y, 8, k/8, nullptr, nullptr);
}
size_t quantize_mxfp4_r8(const float * src, void * dst, int64_t nrows, int64_t n_per_row,
[[maybe_unused]] const float * imatrix,
[[maybe_unused]] const quantize_user_data * user_data) {
GGML_ASSERT(nrows % 8 == 0);
constexpr int kBlockSize = QK_MXFP4;
GGML_ASSERT(n_per_row%kBlockSize == 0);
auto row_size = ggml_row_size(GGML_TYPE_MXFP4, n_per_row);
block_mxfp4_r8 * y = (block_mxfp4_r8 *)dst;
int nblock = n_per_row/QK_MXFP4;
for (int row = 0; row < nrows; row += 8) {
for (int ib = 0; ib < nblock; ++ib) {
memset(&y[ib], 0, sizeof(block_mxfp4_r8));
for (int k = 0; k < 8; ++k) {
const float * xb = src + (row + k)*n_per_row + ib*QK_MXFP4;
float amax = 0;
for (int j = 0; j < QK_MXFP4; ++j) {
float ax = fabsf(xb[j]);
amax = std::max(amax, ax);
}
if (!amax) {
continue;
}
const uint8_t e = (uint8_t) (floorf(log2f(amax)) - 2 + 127);
const float d = GGML_E8M0_TO_FP32_HALF(e);
y[ib].e[k] = e;
for (int j1 = 0; j1 < QK_MXFP4/8; ++j1) {
for (int j2 = 0; j2 < 4; ++j2) {
uint8_t v0 = best_index_mxfp4(d, kvalues_mxfp4, xb[4*j1+j2]);
uint8_t v1 = best_index_mxfp4(d, kvalues_mxfp4, xb[4*j1+j2+QK_MXFP4/2]);
// for each j1 we have 4 values per row with 8 interleaved rows
y[ib].qs[32*j1 + 4*k + j2] = v0 | (v1 << 4);
}
}
}
}
y += nblock;
}
return nrows * row_size;
}
void dequantize_row_mxfp4_r8(const block_mxfp4_r8 * x, float * y, int64_t k) {
constexpr int kBlockSize = QK_MXFP4;
int n_per_row = k/8;
GGML_ASSERT(n_per_row%kBlockSize == 0);
int nblock = k/kBlockSize;
float d[8];
uint32_t aux32[2];
const uint8_t * aux8 = (const uint8_t *)aux32;
float * y8[8];
for (int k = 0; k < 8; ++k) y8[k] = y + k*n_per_row;
for (int ib = 0; ib < nblock; ++ib) {
for (int k = 0; k < 8; ++k) d[k] = GGML_E8M0_TO_FP32_HALF(x[ib].e[k]);
auto qs = (const uint32_t *)x[ib].qs;
for (int j = 0; j < kBlockSize/8; ++j) {
for (int k = 0; k < 8; ++k) {
aux32[0] = qs[8*j+k] & 0x0f0f0f0f;
aux32[1] = (qs[8*j+k] >> 4) & 0x0f0f0f0f;
for (int i = 0; i < 4; ++i) {
y8[k][kBlockSize*ib + 4*j + i ] = d[k] * kvalues_mxfp4[aux8[i+0]];
y8[k][kBlockSize*ib + 4*j + i + kBlockSize/2] = d[k] * kvalues_mxfp4[aux8[i+4]];
}
}
}
}
}
void vec_dot_mxfp4_r8_q8_2_x4(int n, float * s, size_t bs, const void * vx, size_t bx, const void * vy, size_t by, int nrc) {
#if GGML_USE_IQK_MULMAT
if (iqk_mul_mat(1, 1, n, GGML_TYPE_MXFP4_R8, vx, 0, GGML_TYPE_Q8_2_X4, vy, 0, s, 0, 0, 1)) {
return;
}
#endif
GGML_ASSERT(n%QK_MXFP4 == 0);
GGML_ASSERT(nrc == 1);
GGML_UNUSED(bs);
GGML_UNUSED(bx);
GGML_UNUSED(by);
}
namespace {
static void quantize_row_iq4_k_impl_bs128(const int super_block_size, const int block_size,
int n_per_row, const float * x, char * cy,
@ -5276,6 +5347,28 @@ static void modify_q4_0_r8(int64_t k, char * cy) {
}
#endif
static void repack_mxfp4(int nrows, int n_per_row, const block_mxfp4 * x, block_mxfp4_r8 * y, [[maybe_unused]] bool online) {
GGML_ASSERT(nrows%8 == 0);
GGML_ASSERT(n_per_row%QK_MXFP4 == 0);
int nblock = n_per_row/QK_MXFP4;
const block_mxfp4 * x8[8];
for (int row = 0; row < nrows; row += 8) {
for (int k = 0; k < 8; ++k) x8[k] = x + nblock*k;
for (int ib = 0; ib < nblock; ++ib) {
for (int k = 0; k < 8; ++k) {
y[ib].e[k] = x8[k][ib].e;
for (int l = 0; l < 4; ++l) {
for (int i = 0; i < 4; ++i) {
y[ib].qs[32*l+4*k+i] = x8[k][ib].qs[4*l + i];
}
}
}
}
x += 8*nblock;
y += nblock;
}
}
size_t quantize_q4_0_r8(const float * src, void * dst, int64_t nrows, int64_t n_per_row, const float * imatrix,
const quantize_user_data * user_data) {
GGML_ASSERT(nrows%8 == 0);
@ -8421,6 +8514,7 @@ const Repack * get_repack_info(ggml_type type) {
{ GGML_TYPE_Q8_0, { GGML_TYPE_Q8_0_R8, 8, (Repack::repack_func)repack_q8_0} },
{ GGML_TYPE_Q8_K, { GGML_TYPE_Q8_K_R8, 8, (Repack::repack_func)repack_q8_k} },
{ GGML_TYPE_Q8_KV, { GGML_TYPE_Q8_KV_R8, 8, (Repack::repack_func)repack_q8_KV} },
{ GGML_TYPE_MXFP4, { GGML_TYPE_MXFP4_R8, 8, (Repack::repack_func)repack_mxfp4} },
#ifdef __AVX512BF16__
{ GGML_TYPE_BF16, { GGML_TYPE_BF16_R16, 16, (Repack::repack_func)repack_bf16<ggml_bf16_t>}},
{ GGML_TYPE_F16, { GGML_TYPE_BF16_R16, 16, (Repack::repack_func)repack_bf16<ggml_half>} },

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

@ -303,6 +303,12 @@ size_t quantize_q1_0_g128(const float * GGML_RESTRICT src, void * GGML_RESTRICT
void dequantize_row_q1_0_g128(const block_q1_0_g128 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
void 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 quantize_row_mxfp4_r8_ref(const float * GGML_RESTRICT x, block_mxfp4_r8 * GGML_RESTRICT y, int64_t k);
void quantize_row_mxfp4_r8(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
size_t quantize_mxfp4_r8(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix, const struct quantize_user_data * use_data);
void dequantize_row_mxfp4_r8(const block_mxfp4_r8 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
void vec_dot_mxfp4_r8_q8_2_x4(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 iqk_quantize_row_q8_K(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k);
void quantize_row_q8_K64_ref(const float * GGML_RESTRICT x, block_q8_K64 * GGML_RESTRICT y, int64_t k);
void quantize_row_q8_K64(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);

1
include/llama.h
View File

@ -241,6 +241,7 @@ extern "C" {
LLAMA_FTYPE_MOSTLY_IQ5_K_R4 = 341, // except 1d tensors
LLAMA_FTYPE_MOSTLY_IQ4_KS_R4 = 345, // except 1d tensors
LLAMA_FTYPE_MOSTLY_IQ5_KS_R4 = 350, // except 1d tensors
LLAMA_FTYPE_MOSTLY_MXFP4_R8 = 351, // except 1d tensors
LLAMA_FTYPE_MOSTLY_Q8_KV_R8 = 398, // except 1d tensors
LLAMA_FTYPE_MOSTLY_Q8_K_R8 = 399, // except 1d tensors

2
src/llama-quantize.cpp
View File

@ -110,6 +110,7 @@ std::pair<ggml_type, int> interleaved_properties(ggml_type type) {
{ GGML_TYPE_IQ4_KS_R4, { GGML_TYPE_IQ4_KS, 4} },
{ GGML_TYPE_IQ5_KS_R4, { GGML_TYPE_IQ5_KS, 4} },
{ GGML_TYPE_IQ5_K_R4, { GGML_TYPE_IQ5_K, 4} },
{ GGML_TYPE_MXFP4_R8, { GGML_TYPE_MXFP4_R8, 8} },
{ GGML_TYPE_Q8_KV_R8, { GGML_TYPE_Q8_KV, 8} },
{ GGML_TYPE_Q8_K_R8, { GGML_TYPE_Q8_0, 8} },
{ GGML_TYPE_BF16_R16, { GGML_TYPE_BF16, 16} },
@ -1089,6 +1090,7 @@ static void llama_model_quantize_internal(const std::string & fname_inp, const s
case LLAMA_FTYPE_MOSTLY_Q6_0_R4: default_type = GGML_TYPE_Q6_0_R4; break;
case LLAMA_FTYPE_MOSTLY_Q8_0_R8: default_type = GGML_TYPE_Q8_0_R8; break;
case LLAMA_FTYPE_MOSTLY_MXFP4: default_type = GGML_TYPE_MXFP4; break;
case LLAMA_FTYPE_MOSTLY_MXFP4_R8:default_type = GGML_TYPE_MXFP4_R8;break;
case LLAMA_FTYPE_MOSTLY_Q1_0_G128: default_type = GGML_TYPE_Q1_0_G128; break;
case LLAMA_FTYPE_MOSTLY_IQ4_XS: default_type = GGML_TYPE_IQ4_XS; break;
case LLAMA_FTYPE_MOSTLY_IQ4_KS: default_type = GGML_TYPE_IQ4_KS; break;