diff --git a/examples/quantize/quantize.cpp b/examples/quantize/quantize.cpp index 7ff87e18..e5e62314 100644 --- a/examples/quantize/quantize.cpp +++ b/examples/quantize/quantize.cpp @@ -29,6 +29,7 @@ static const std::vector 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", }, diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h index fe9d09ca..d16a2d02 100644 --- a/ggml/include/ggml.h +++ b/ggml/include/ggml.h @@ -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 diff --git a/ggml/src/ggml-common.h b/ggml/src/ggml-common.h index 708f46f2..f018d58f 100644 --- a/ggml/src/ggml-common.h +++ b/ggml/src/ggml-common.h @@ -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 diff --git a/ggml/src/ggml-quants.c b/ggml/src/ggml-quants.c index 7533985a..7d7acbe1 100644 --- a/ggml/src/ggml-quants.c +++ b/ggml/src/ggml-quants.c @@ -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; diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index ee2e3167..c21d19a7 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -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; diff --git a/ggml/src/iqk/iqk_gemm_legacy_quants.cpp b/ggml/src/iqk/iqk_gemm_legacy_quants.cpp index adbe8914..0961a701 100644 --- a/ggml/src/iqk/iqk_gemm_legacy_quants.cpp +++ b/ggml/src/iqk/iqk_gemm_legacy_quants.cpp @@ -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 +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 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 +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 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 +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(n, vx, bx, info, nrc_x); +} +#endif + template 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; #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; diff --git a/ggml/src/iqk/iqk_mul_mat.cpp b/ggml/src/iqk/iqk_mul_mat.cpp index 21f68e2d..14ed1634 100644 --- a/ggml/src/iqk/iqk_mul_mat.cpp +++ b/ggml/src/iqk/iqk_mul_mat.cpp @@ -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: diff --git a/ggml/src/iqk/iqk_quantize.cpp b/ggml/src/iqk/iqk_quantize.cpp index 35b21c2e..015205c9 100644 --- a/ggml/src/iqk/iqk_quantize.cpp +++ b/ggml/src/iqk/iqk_quantize.cpp @@ -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_TYPE_F16, { GGML_TYPE_BF16_R16, 16, (Repack::repack_func)repack_bf16} }, diff --git a/ggml/src/iqk/iqk_quantize.h b/ggml/src/iqk/iqk_quantize.h index b5fb8a31..93ce1e1c 100644 --- a/ggml/src/iqk/iqk_quantize.h +++ b/ggml/src/iqk/iqk_quantize.h @@ -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); diff --git a/include/llama.h b/include/llama.h index 4573775d..ed744d29 100644 --- a/include/llama.h +++ b/include/llama.h @@ -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 diff --git a/src/llama-quantize.cpp b/src/llama-quantize.cpp index 0c12e40e..39a597b6 100644 --- a/src/llama-quantize.cpp +++ b/src/llama-quantize.cpp @@ -110,6 +110,7 @@ std::pair 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;