From ce77728c06fa6b5898f8ef557768c2fe3e129d0a Mon Sep 17 00:00:00 2001 From: Georgi Gerganov Date: Fri, 21 Aug 2026 14:02:03 +0300 Subject: [PATCH] Revert "sycl : add Q2_K reordered MMVQ and ESIMD kernels (llama/26336)" (llama/27486) This reverts commit ff14356e0caf6988f61f1f15f9dfe7d5ab398271. --- ggml/src/ggml-sycl/convert.cpp | 25 +-------- ggml/src/ggml-sycl/dequantize.hpp | 41 --------------- ggml/src/ggml-sycl/dmmv.cpp | 27 +--------- ggml/src/ggml-sycl/esimd.hpp | 87 ------------------------------- ggml/src/ggml-sycl/ggml-sycl.cpp | 2 - ggml/src/ggml-sycl/mmvq.cpp | 75 +------------------------- ggml/src/ggml-sycl/quants.hpp | 23 -------- ggml/src/ggml-sycl/vecdotq.hpp | 33 ------------ 8 files changed, 4 insertions(+), 309 deletions(-) diff --git a/ggml/src/ggml-sycl/convert.cpp b/ggml/src/ggml-sycl/convert.cpp index b660b56ab..9ec927695 100644 --- a/ggml/src/ggml-sycl/convert.cpp +++ b/ggml/src/ggml-sycl/convert.cpp @@ -76,19 +76,6 @@ static void dequantize_row_q2_K_sycl(const void *vx, dst_t *y, const int64_t k, #endif } -template -static void dequantize_row_q2_K_sycl_reorder(const void *vx, dst_t *y, const int64_t k, - dpct::queue_ptr stream) { - const int64_t nb = k / QK_K; - - dpct::has_capability_or_fail(stream->get_device(), { sycl::aspect::fp16 }); - stream->parallel_for( - sycl::nd_range<3>(sycl::range<3>(1, 1, nb) * sycl::range<3>(1, 1, 64), sycl::range<3>(1, 1, 64)), - [=](sycl::nd_item<3> item_ct1) { - dequantize_block_q2_K_reorder(vx, y, item_ct1, nb); - }); -} - template static void dequantize_row_q3_K_sycl(const void *vx, dst_t *y, const int64_t k, dpct::queue_ptr stream) { @@ -680,11 +667,7 @@ to_fp16_sycl_t ggml_get_to_fp16_sycl(ggml_type type, ggml_tensor * dst) { return dequantize_block_sycl; } case GGML_TYPE_Q2_K: - if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { - return dequantize_row_q2_K_sycl_reorder; - } else { - return dequantize_row_q2_K_sycl; - } + return dequantize_row_q2_K_sycl; case GGML_TYPE_Q3_K: if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { return dequantize_row_q3_K_sycl_reorder; @@ -770,11 +753,7 @@ to_fp32_sycl_t ggml_get_to_fp32_sycl(ggml_type type, ggml_tensor *dst) { return dequantize_block_sycl; } case GGML_TYPE_Q2_K: - if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { - return dequantize_row_q2_K_sycl_reorder; - } else { - return dequantize_row_q2_K_sycl; - } + return dequantize_row_q2_K_sycl; case GGML_TYPE_Q3_K: if (dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { return dequantize_row_q3_K_sycl_reorder; diff --git a/ggml/src/ggml-sycl/dequantize.hpp b/ggml/src/ggml-sycl/dequantize.hpp index 1b13e0f1a..876ba1b44 100644 --- a/ggml/src/ggml-sycl/dequantize.hpp +++ b/ggml/src/ggml-sycl/dequantize.hpp @@ -943,47 +943,6 @@ static void dequantize_block_q2_K(const void * __restrict__ vx, dst_t * __restri } -template -static void dequantize_block_q2_K_reorder(const void * __restrict__ vx, dst_t * __restrict__ yy, - const sycl::nd_item<3> & item_ct1, int64_t n_blocks) { -#if QK_K == 256 - const int64_t i = item_ct1.get_group(2); - if (i >= n_blocks) { - return; - } - - const uint8_t * base = static_cast(vx); - const size_t qs_offset = i * (QK_K / 4); - const size_t scales_offset = n_blocks * (QK_K / 4) + i * (QK_K / 16); - const size_t dm_offset = n_blocks * (QK_K / 4) + n_blocks * (QK_K / 16) + i * sizeof(ggml_half2); - - const uint8_t * qs = base + qs_offset; - const uint8_t * scales = base + scales_offset; - const ggml_half2 * dm = reinterpret_cast(base + dm_offset); - - const int64_t tid = item_ct1.get_local_id(2); - const int64_t n = tid / 32; - const int64_t l = tid - 32 * n; - const int64_t is = 8 * n + l / 16; - - const uint8_t q = qs[32 * n + l]; - dst_t * y = yy + i * QK_K + 128 * n; - - const float dall = (*dm)[0]; - const float dmin = (*dm)[1]; - y[l+ 0] = dall * (scales[is+0] & 0xF) * ((q >> 0) & 3) - dmin * (scales[is+0] >> 4); - y[l+32] = dall * (scales[is+2] & 0xF) * ((q >> 2) & 3) - dmin * (scales[is+2] >> 4); - y[l+64] = dall * (scales[is+4] & 0xF) * ((q >> 4) & 3) - dmin * (scales[is+4] >> 4); - y[l+96] = dall * (scales[is+6] & 0xF) * ((q >> 6) & 3) - dmin * (scales[is+6] >> 4); -#else - GGML_UNUSED(vx); - GGML_UNUSED(yy); - GGML_UNUSED(item_ct1); - GGML_UNUSED(n_blocks); - GGML_ABORT("Q2_K reorder dequantize not supported for QK_K != 256"); -#endif -} - template static void dequantize_block_q3_K(const void * __restrict__ vx, dst_t * __restrict__ yy, const sycl::nd_item<3> &item_ct1) { diff --git a/ggml/src/ggml-sycl/dmmv.cpp b/ggml/src/ggml-sycl/dmmv.cpp index d47d6831a..fdcadbf91 100644 --- a/ggml/src/ggml-sycl/dmmv.cpp +++ b/ggml/src/ggml-sycl/dmmv.cpp @@ -1921,23 +1921,6 @@ ESIMD_INLINE void dequantize_mul_mat_vec_reorder_esimd( } } -static void dequantize_mul_mat_vec_q2_K_sycl_reorder_esimd(const void *vx, const float *y, - float *dst, const int ncols, - const int nrows, - dpct::queue_ptr stream) { - GGML_ASSERT(ncols % QK_K == 0); - const int workgroups = (nrows + 1) / 2; - stream->submit([&](sycl::handler &h) { - sycl::local_accessor lmem(sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE * 2), h); - h.parallel_for( - sycl::nd_range<1>(sycl::range<1>((size_t)workgroups * GGML_SYCL_DMMV_ESIMD_WG_SIZE), sycl::range<1>(GGML_SYCL_DMMV_ESIMD_WG_SIZE)), - [=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] { - dequantize_mul_mat_vec_reorder_esimd( - vx, y, dst, ncols, nrows, lmem, it); - }); - }); -} - static void dequantize_mul_mat_vec_q3_K_sycl_reorder_esimd(const void *vx, const float *y, float *dst, const int ncols, const int nrows, @@ -2128,15 +2111,7 @@ void ggml_sycl_op_dequantize_mul_mat_vec( case GGML_TYPE_Q2_K: if ((ggml_tensor_extra_gpu *) dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { -#ifdef GGML_SYCL_DMMV_HAS_ESIMD - if (g_ggml_sycl_enable_esimd) { - dequantize_mul_mat_vec_q2_K_sycl_reorder_esimd(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); - } - else -#endif - { - dequantize_mul_mat_vec_q2_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); - } + dequantize_mul_mat_vec_q2_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); } else { dequantize_mul_mat_vec_q2_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); } diff --git a/ggml/src/ggml-sycl/esimd.hpp b/ggml/src/ggml-sycl/esimd.hpp index 0485ff0ce..04e596ed3 100644 --- a/ggml/src/ggml-sycl/esimd.hpp +++ b/ggml/src/ggml-sycl/esimd.hpp @@ -61,93 +61,6 @@ static ESIMD_INLINE void unpack_scale_min_k4( min_f = convert(m) * (-dmin); } -// --------------------------------------------------------------------------- -// Q2_K, SOA reorder layout produced by reorder_qw_q2_k: -// [qs: nb*(QK_K/4)] [scales: nb*(QK_K/16)] [dm: nb*sizeof(half2)] -// with nb = nrows*num_blocks_per_row. -// -// 2 bits per weight. The 8 output chunks of 32 (matching dequantize_row_q2_K) -// map to super-chunk s (0..7): byte base 32*(s/4) into the 64-byte qs array, -// bit shift 2*(s%4); the low 16 lanes use scales[2s], the high 16 use -// scales[2s+1], with dl = d*(sc & 0xF), ml = dmin*(sc >> 4), deq = dl*q - ml. -// --------------------------------------------------------------------------- -template <> struct esimd_reorder_q_traits { - struct ptrs { - const uint8_t * qs; - const uint8_t * scales; - const sycl::half * dm; - }; - - static ESIMD_INLINE ptrs make_ptrs(const void * vx, size_t nb) { - const uint8_t * qs = (const uint8_t *) vx; - const uint8_t * scales = qs + nb * (QK_K / 4); - const sycl::half * dm = (const sycl::half *) (scales + nb * (QK_K / 16)); - return { qs, scales, dm }; - } - - static ESIMD_INLINE void mac_pair( - const ptrs & pa, size_t bia, - const ptrs & pb, size_t bib, bool has_b, - sycl::ext::intel::esimd::simd & y_vec, - sycl::ext::intel::esimd::simd & acc_a, - sycl::ext::intel::esimd::simd & acc_b) { - using namespace sycl::ext::intel::esimd; - - simd qs_a = block_load(pa.qs + bia * (QK_K / 4)); - simd qs_b = 0; - simd scales_a = block_load(pa.scales + bia * (QK_K / 16)); - simd scales_b = 0; - - const float dall_a = (float) pa.dm[bia * 2 + 0]; - const float dmin_a = (float) pa.dm[bia * 2 + 1]; - float dall_b = 0.0f; - float dmin_b = 0.0f; - if (has_b) { - qs_b = block_load(pb.qs + bib * (QK_K / 4)); - scales_b = block_load(pb.scales + bib * (QK_K / 16)); - dall_b = (float) pb.dm[bib * 2 + 0]; - dmin_b = (float) pb.dm[bib * 2 + 1]; - } - - // per-chunk scale (d * (sc & 0xF)) and min (-dmin * (sc >> 4)), all 16 codes; - // min carries the negation so the dequant epilogue adds (matches Q4_K/Q5_K) - simd scale_f_a = convert(scales_a & simd(0x0F)) * dall_a; - simd min_f_a = convert(scales_a >> simd(4)) * (-dmin_a); - simd scale_f_b = convert(scales_b & simd(0x0F)) * dall_b; - simd min_f_b = convert(scales_b >> simd(4)) * (-dmin_b); - -#pragma unroll - for (int s = 0; s < 8; ++s) { - const int byte_base = 32 * (s / 4); - const uint8_t shift = (uint8_t) (2 * (s % 4)); - simd y_s = y_vec.select<32, 1>(s * 32); - - simd qa = (qs_a.select<32, 1>(byte_base) >> shift) & simd(3); - simd qb = (qs_b.select<32, 1>(byte_base) >> shift) & simd(3); - - const float scale_a_lo = scale_f_a[2 * s + 0]; - const float scale_a_hi = scale_f_a[2 * s + 1]; - const float min_a_lo = min_f_a[2 * s + 0]; - const float min_a_hi = min_f_a[2 * s + 1]; - const float scale_b_lo = scale_f_b[2 * s + 0]; - const float scale_b_hi = scale_f_b[2 * s + 1]; - const float min_b_lo = min_f_b[2 * s + 0]; - const float min_b_hi = min_f_b[2 * s + 1]; - - simd scale_vec_a = splat_lo_hi(scale_a_lo, scale_a_hi); - simd min_vec_a = splat_lo_hi(min_a_lo, min_a_hi); - simd scale_vec_b = splat_lo_hi(scale_b_lo, scale_b_hi); - simd min_vec_b = splat_lo_hi(min_b_lo, min_b_hi); - - simd deq_a = convert(qa) * scale_vec_a + min_vec_a; - simd deq_b = convert(qb) * scale_vec_b + min_vec_b; - - acc_a += y_s * deq_a; - acc_b += y_s * deq_b; - } - } -}; - // --------------------------------------------------------------------------- // Q3_K, SOA reorder layout produced by reorder_qw_q3_k: // [qs: nb*(QK_K/4)] [hmask: nb*(QK_K/8)] [scales: nb*12] [d: nb*sizeof(half)] diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 3f82020f4..de56ea5b9 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -3796,7 +3796,6 @@ inline bool ggml_sycl_supports_reorder_mmvq(enum ggml_type type) { case GGML_TYPE_Q1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q8_0: - case GGML_TYPE_Q2_K: case GGML_TYPE_Q3_K: case GGML_TYPE_Q4_K: case GGML_TYPE_Q5_K: @@ -3810,7 +3809,6 @@ inline bool ggml_sycl_supports_reorder_mmvq(enum ggml_type type) { static bool ggml_sycl_supports_reorder_esimd(enum ggml_type type) { #ifdef GGML_SYCL_DMMV_HAS_ESIMD switch (type) { - case GGML_TYPE_Q2_K: case GGML_TYPE_Q3_K: case GGML_TYPE_Q4_K: case GGML_TYPE_Q5_K: diff --git a/ggml/src/ggml-sycl/mmvq.cpp b/ggml/src/ggml-sycl/mmvq.cpp index bfccb4b08..123b2a2f0 100644 --- a/ggml/src/ggml-sycl/mmvq.cpp +++ b/ggml/src/ggml-sycl/mmvq.cpp @@ -1401,64 +1401,6 @@ static void mul_mat_vec_q2_K_q8_1_sycl_switch_ncols( } } -static void reorder_mul_mat_vec_q2_k_q8_1_sycl(const void * vx, const void * vy, float * dst, const int ncols, - const int nrows, dpct::queue_ptr stream) { - GGML_ASSERT(ncols % QK_K == 0); - - // Round up to a whole number of subgroup-sized workgroups; out-of-range rows are skipped inside the kernel. - constexpr size_t num_subgroups = WARP_SIZE; - const int block_num_y = ceil_div(nrows, GGML_SYCL_MMV_Y * (int) num_subgroups); - const sycl::range<3> block_nums(1, 1, block_num_y); - const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, num_subgroups * WARP_SIZE); - - stream->submit([&](sycl::handler & cgh) { - cgh.parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), - [=](sycl::nd_item<3> nd_item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] { - mul_mat_vec_q_reorder>(vx, vy, dst, ncols, nrows, - nd_item); - }); - }); -} - -template -static void reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols( - const void * vx, const void * vy, float * dst, - const int ncols, const int nrows, - const int stride_col_y_bytes, const int stride_col_dst, - dpct::queue_ptr stream) { - GGML_ASSERT(ncols % QK_K == 0); - constexpr size_t num_subgroups = WARP_SIZE; - const int block_num_y = ceil_div(nrows, GGML_SYCL_MMV_Y * (int) num_subgroups); - const sycl::range<3> block_nums(1, 1, block_num_y); - const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, num_subgroups * WARP_SIZE); - - stream->submit([&](sycl::handler & cgh) { - cgh.parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), - [=](sycl::nd_item<3> nd_item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] { - mul_mat_vec_q_reorder_ncols, ncols_dst>( - vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, nd_item); - }); - }); -} - -static void reorder_mul_mat_vec_q2_k_q8_1_sycl_switch_ncols( - const void * vx, const void * vy, float * dst, - const int ncols, const int nrows, const int ncols_dst, - const int stride_col_y_bytes, const int stride_col_dst, - dpct::queue_ptr stream) { - switch (ncols_dst) { - case 1: reorder_mul_mat_vec_q2_k_q8_1_sycl(vx, vy, dst, ncols, nrows, stream); break; - case 2: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<2>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; - case 3: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<3>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; - case 4: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<4>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; - case 5: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<5>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; - case 6: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<6>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; - case 7: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<7>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; - case 8: reorder_mul_mat_vec_q2_k_q8_1_sycl_ncols<8>(vx, vy, dst, ncols, nrows, stride_col_y_bytes, stride_col_dst, stream); break; - default: GGML_ABORT("unsupported ncols_dst=%d for Q2_K reorder multi-col MMVQ", ncols_dst); - } -} - static void mul_mat_vec_q3_K_q8_1_sycl(const void *vx, const void *vy, float *dst, const int ncols, const int nrows, @@ -2355,21 +2297,7 @@ void ggml_sycl_op_mul_mat_vec_q(ggml_backend_sycl_context & ctx, const ggml_tens } break; case GGML_TYPE_Q2_K: - if ((ggml_tensor_extra_gpu *) dst->src[0]->extra && - ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { - if (i == 0 && src1_ncols > 1 && src1_ncols <= 8) { - const int stride_col_y_bytes = src1_padded_col_size * q8_1_ts / q8_1_bs; - const int stride_col_dst = dst->ne[0]; - GGML_SYCL_DEBUG("Calling reorder_mul_mat_vec_q2_k_q8_1_sycl_switch_ncols ncols=%d\n", (int)src1_ncols); - reorder_mul_mat_vec_q2_k_q8_1_sycl_switch_ncols( - src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, - src1_ncols, stride_col_y_bytes, stride_col_dst, stream); - return; - } else { - GGML_SYCL_DEBUG("Calling reorder_mul_mat_vec_q2_k_q8_1_sycl\n"); - reorder_mul_mat_vec_q2_k_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream); - } - } else if (i == 0 && src1_ncols > 1 && src1_ncols <= 8) { + if (i == 0 && src1_ncols > 1 && src1_ncols <= 8) { const int stride_col_y = src1_padded_col_size / QK8_1; const int stride_col_dst = dst->ne[0]; GGML_SYCL_DEBUG("Calling mul_mat_vec_q2_K_q8_1_sycl_switch_ncols ncols=%d\n", (int)src1_ncols); @@ -2378,7 +2306,6 @@ void ggml_sycl_op_mul_mat_vec_q(ggml_backend_sycl_context & ctx, const ggml_tens src1_ncols, stride_col_y, stride_col_dst, stream); return; } else if (i == 0 || src1_ncols == 1) { - GGML_SYCL_DEBUG("Calling mul_mat_vec_q2_K_q8_1_sycl\n"); mul_mat_vec_q2_K_q8_1_sycl(src0_dd_i, src1_ddq_i_bs, dst_dd_i_bs, ne00, row_diff, stream); } break; diff --git a/ggml/src/ggml-sycl/quants.hpp b/ggml/src/ggml-sycl/quants.hpp index a26a6ce6e..95287f175 100644 --- a/ggml/src/ggml-sycl/quants.hpp +++ b/ggml/src/ggml-sycl/quants.hpp @@ -58,29 +58,6 @@ template <> struct block_q_t { static constexpr int block_to_q8_1_ratio() { return traits::qk / QK8_1; } }; -template <> struct block_q_t { - struct traits { - static constexpr uint32_t qk = QK_K; - static constexpr uint32_t qi = QI2_K; - static constexpr uint32_t qr = QR2_K; - static constexpr uint32_t vdr_mmvq = 1; - }; - - // Reordered layout: [qs (QK_K/4 per block)] [scales (QK_K/16 per block)] [dm] - static constexpr std::pair get_block_offset(const int block_index, const int /* n_blocks */) { - return { block_index * (QK_K / 4), 0 }; - } - - static constexpr std::pair get_d_offset(int nrows, int ncols, const int block_index) { - auto nblocks = (nrows * (ncols / QK_K)); - auto total_qs_bytes = nblocks * (QK_K / 4); - return { total_qs_bytes + block_index * (QK_K / 16), - total_qs_bytes + nblocks * (QK_K / 16) + block_index * sizeof(ggml_half2) }; - } - - static constexpr int block_to_q8_1_ratio() { return traits::qk / QK8_1; } -}; - template <> struct block_q_t { struct traits { static constexpr uint32_t qk = QK_K; diff --git a/ggml/src/ggml-sycl/vecdotq.hpp b/ggml/src/ggml-sycl/vecdotq.hpp index 3ad4cee93..c11a6e8f9 100644 --- a/ggml/src/ggml-sycl/vecdotq.hpp +++ b/ggml/src/ggml-sycl/vecdotq.hpp @@ -429,39 +429,6 @@ template <> struct reorder_vec_dot_q_sycl { } }; -template <> struct reorder_vec_dot_q_sycl { - static constexpr ggml_type gtype = GGML_TYPE_Q2_K; - - using q2_k_block = ggml_sycl_reordered::block_q_t; - using q2_k_traits = typename q2_k_block::traits; - - __dpct_inline__ float operator()(const void * __restrict__ vbq, const std::pair ibx_offset, - const std::pair d_offset, const int8_t * q8_1_quant_ptr, - const sycl::half2 * q8_1_ds, const int & iqs) { - const uint8_t * base = static_cast(vbq); - const uint8_t * qs = base + ibx_offset.first; - const uint8_t * scales = base + d_offset.first; - const ggml_half2 * dm = reinterpret_cast(base + d_offset.second); - - const int bq8_offset = QR2_K * (iqs / QI8_1); - const int scale_offset = iqs - iqs % QI8_1 + (iqs % QI8_1) / (QI8_1 / 2); - - const int v = get_int_from_uint8_aligned(qs, iqs); - - int u[QR2_K]; - float d8[QR2_K]; - -#pragma unroll - for (int i = 0; i < QR2_K; ++i) { - const int8_t * quant_base_ptr = q8_1_quant_ptr + (bq8_offset + i) * QK8_1; - u[i] = get_int_from_int8_aligned(quant_base_ptr, iqs % QI8_1); - d8[i] = (*(q8_1_ds + bq8_offset + i))[0]; - } - - return vec_dot_q2_K_q8_1_impl_mmvq(v, u, scales + scale_offset, *dm, d8); - } -}; - template <> struct reorder_vec_dot_q_sycl { static constexpr ggml_type gtype = GGML_TYPE_Q3_K;