#include <sycl/ext/oneapi/bfloat16.hpp>
#define GGML_SYCL_DMMV_HAS_BF16
#endif
+ #include <sycl/ext/intel/esimd.hpp>
+ #include "esimd.hpp"
+ #define GGML_SYCL_DMMV_HAS_ESIMD
#endif
static void convert_f16(const void * vx, const int64_t ib, const int iqs, dfloat2 & v){
});
}
+#ifdef GGML_SYCL_DMMV_HAS_ESIMD
+using ggml_sycl_esimd::GGML_SYCL_DMMV_ESIMD_WG_SIZE;
+
+// generic reordered dequantize-matvec: each work-group owns a pair of
+// consecutive output rows and updates one 32-wide accumulator per row
+template <ggml_type T>
+ESIMD_INLINE void dequantize_mul_mat_vec_reorder_esimd(
+ const void * vx, const float * y, float * dst,
+ const int ncols, const int nrows,
+ sycl::local_accessor<float, 1> lmem,
+ const sycl::nd_item<1> & it) {
+ using namespace sycl::ext::intel::esimd;
+ using traits = ggml_sycl_esimd::esimd_reorder_q_traits<T>;
+
+ const int num_blocks_per_row = ncols / QK_K;
+ const size_t nb = (size_t) nrows * num_blocks_per_row;
+ const auto ps = traits::make_ptrs(vx, nb);
+
+ const int tid = it.get_local_id(0);
+ const int row_pair = it.get_group(0);
+ const int row0 = row_pair * 2; // two consecutive output rows
+ const bool has_row1 = row0 + 1 < nrows;
+
+ // one 32-wide accumulator per output row (small footprint, no spill)
+ simd<float, 32> acc0 = 0.0f;
+ simd<float, 32> acc1 = 0.0f;
+
+ for (int ib = tid; ib < num_blocks_per_row; ib += GGML_SYCL_DMMV_ESIMD_WG_SIZE) {
+ simd<float, 256> y_vec = block_load<float, 256>(y + (size_t) ib * QK_K);
+
+ const size_t bi0 = (size_t) (row0 + 0) * num_blocks_per_row + ib;
+ const size_t bi1 = (size_t) (row0 + 1) * num_blocks_per_row + ib;
+
+ traits::mac_pair(ps, bi0, ps, bi1, has_row1, y_vec, acc0, acc1);
+ }
+
+ lmem[tid * 2 + 0] = reduce<float>(acc0, std::plus<>{});
+ lmem[tid * 2 + 1] = reduce<float>(acc1, std::plus<>{});
+ it.barrier(sycl::access::fence_space::local_space);
+
+ if (tid == 0) {
+ float sum0 = 0.0f;
+ float sum1 = 0.0f;
+ for (int p = 0; p < GGML_SYCL_DMMV_ESIMD_WG_SIZE; ++p) {
+ sum0 += lmem[p * 2 + 0];
+ sum1 += lmem[p * 2 + 1];
+ }
+ dst[row0 + 0] = sum0;
+ if (has_row1) {
+ dst[row0 + 1] = sum1;
+ }
+ }
+}
+
+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,
+ dpct::queue_ptr stream) {
+ GGML_ASSERT(ncols % QK_K == 0);
+ const int workgroups = (nrows + 1) / 2;
+ stream->submit([&](sycl::handler &h) {
+ sycl::local_accessor<float, 1> 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<GGML_TYPE_Q3_K>(
+ vx, y, dst, ncols, nrows, lmem, it);
+ });
+ });
+}
+
+static void dequantize_mul_mat_vec_q4_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<float, 1> 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<GGML_TYPE_Q4_K>(
+ vx, y, dst, ncols, nrows, lmem, it);
+ });
+ });
+}
+
+static void dequantize_mul_mat_vec_q6_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<float, 1> 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<GGML_TYPE_Q6_K>(
+ vx, y, dst, ncols, nrows, lmem, it);
+ });
+ });
+}
+
+#endif // GGML_SYCL_DMMV_HAS_ESIMD
+
static void dequantize_mul_mat_vec_q4_K_sycl_reorder(const void *vx, const float *y,
float *dst, const int ncols,
const int nrows,
case GGML_TYPE_Q3_K:
if ((ggml_tensor_extra_gpu *) dst->src[0]->extra &&
((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
- dequantize_mul_mat_vec_q3_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+#ifdef GGML_SYCL_DMMV_HAS_ESIMD
+ if (g_ggml_sycl_enable_esimd) {
+ dequantize_mul_mat_vec_q3_K_sycl_reorder_esimd(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+ }
+ else
+#endif
+ {
+ dequantize_mul_mat_vec_q3_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+ }
} else {
dequantize_mul_mat_vec_q3_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
}
case GGML_TYPE_Q4_K:
if ((ggml_tensor_extra_gpu *) dst->src[0]->extra &&
((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
- dequantize_mul_mat_vec_q4_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+#ifdef GGML_SYCL_DMMV_HAS_ESIMD
+ if (g_ggml_sycl_enable_esimd) {
+ dequantize_mul_mat_vec_q4_K_sycl_reorder_esimd(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+ }
+ else
+#endif
+ {
+ dequantize_mul_mat_vec_q4_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+ }
} else {
dequantize_mul_mat_vec_q4_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
}
case GGML_TYPE_Q6_K:
if ((ggml_tensor_extra_gpu *) dst->src[0]->extra &&
((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) {
- dequantize_mul_mat_vec_q6_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+#ifdef GGML_SYCL_DMMV_HAS_ESIMD
+ if (g_ggml_sycl_enable_esimd) {
+ dequantize_mul_mat_vec_q6_K_sycl_reorder_esimd(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+ }
+ else
+#endif
+ {
+ dequantize_mul_mat_vec_q6_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+ }
} else {
dequantize_mul_mat_vec_q6_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
}
--- /dev/null
+//
+// MIT license
+// Copyright (C) 2026 Intel Corporation
+// SPDX-License-Identifier: MIT
+//
+
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+
+#ifndef GGML_SYCL_ESIMD_HPP
+#define GGML_SYCL_ESIMD_HPP
+
+#include <sycl/ext/intel/esimd.hpp>
+
+#include "common.hpp"
+
+namespace ggml_sycl_esimd {
+
+constexpr int GGML_SYCL_DMMV_ESIMD_WG_SIZE = 4;
+
+//
+// Shared ESIMD building blocks for the reordered K-quant dequantize-matvec
+// kernels.
+//
+// The reordered K-quant ESIMD matvec kernels share one skeleton: per super-block,
+// load a 256-float activation slice, load one weight block, dequantize it into 8
+// chunks of 32 and MAC each chunk against the matching activation slice, then
+// reduce and run a lane-0 epilogue.
+//
+// Each K-quant kernel emits exactly 8 chunks of 32 mapping to activation slices
+// 0..7, so the per-block work is captured by esimd_reorder_q_traits<T>::mac_pair,
+// which dequantizes two weight blocks and MACs both against a shared activation
+// vector with the two FMA chains interleaved (co-scheduled to hide FMA latency).
+// The "pair" is the (row0,row1) row pair owned by one work-group, so the
+// layout+dequant is written once per quant type here.
+//
+
+template <ggml_type T> struct esimd_reorder_q_traits;
+
+// build a 32-lane vector whose low 16 lanes are `lo` and high 16 are `hi`
+// (a super-chunk splits into two 16-wide halves with distinct scale/min codes).
+static ESIMD_INLINE sycl::ext::intel::esimd::simd<float, 32> splat_lo_hi(float lo, float hi) {
+ using namespace sycl::ext::intel::esimd;
+ simd<float, 32> v;
+ v.select<16, 1>(0) = lo;
+ v.select<16, 1>(16) = hi;
+ return v;
+}
+
+// unpack one block of Q4_K/Q5_K scale/min codes (get_scale_min_k4 layout) into 8
+// float scales (dall * sc) and 8 float mins (-dmin * m); the min carries the
+// negation so the dequant epilogue adds.
+static ESIMD_INLINE void unpack_scale_min_k4(
+ sycl::ext::intel::esimd::simd<uint8_t, 12> scales, float dall, float dmin,
+ sycl::ext::intel::esimd::simd<float, 8> & scale_f,
+ sycl::ext::intel::esimd::simd<float, 8> & min_f) {
+ using namespace sycl::ext::intel::esimd;
+ simd<uint8_t, 8> sc = 0;
+ simd<uint8_t, 8> m = 0;
+ simd<uint8_t, 4> scale_lo = scales.select<4, 1>(0);
+ simd<uint8_t, 4> min_lo = scales.select<4, 1>(4);
+ simd<uint8_t, 4> hi_bits = scales.select<4, 1>(8);
+ sc.select<4, 1>(0) = scale_lo & simd<uint8_t, 4>(0x3F);
+ sc.select<4, 1>(4) = (hi_bits & simd<uint8_t, 4>(0x0F)) |
+ ((scale_lo >> simd<uint8_t, 4>(6)) << simd<uint8_t, 4>(4));
+ m.select<4, 1>(0) = min_lo & simd<uint8_t, 4>(0x3F);
+ m.select<4, 1>(4) = (hi_bits >> simd<uint8_t, 4>(4)) |
+ ((min_lo >> simd<uint8_t, 4>(6)) << simd<uint8_t, 4>(4));
+ scale_f = convert<float>(sc) * dall;
+ min_f = convert<float>(m) * (-dmin);
+}
+
+// ---------------------------------------------------------------------------
+// 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)]
+// with nb = nrows*num_blocks_per_row. Single super-block scale d, no dmin.
+//
+// 3 bits per weight: 2 low bits in qs, 1 high bit in hmask. The 8 output chunks
+// of 32 (matching dequantize_row_q3_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
+// scale code 2s, the high 16 use 2s+1. hmask is a 32-byte array (like Q5_K's
+// qh) where chunk s uses bit s of the same 32 bytes, but INVERTED: the value is
+// (q & 3) - (hmask_bit_set ? 0 : 4), i.e. (q & 3) + 4*bit - 4.
+//
+// The 16 6-bit scale codes are packed into 12 bytes (get_scale_min layout for
+// Q3_K): low nibbles from bytes 0..7, high 2 bits from bytes 8..11 shifted by
+// 0/2/4/6; the dequant scale is d * (code - 32).
+// ---------------------------------------------------------------------------
+template <> struct esimd_reorder_q_traits<GGML_TYPE_Q3_K> {
+ struct ptrs {
+ const uint8_t * qs;
+ const uint8_t * hmask;
+ const uint8_t * scales;
+ const sycl::half * d;
+ };
+
+ static ESIMD_INLINE ptrs make_ptrs(const void * vx, size_t nb) {
+ const uint8_t * qs = (const uint8_t *) vx;
+ const uint8_t * hmask = qs + nb * (QK_K / 4);
+ const uint8_t * scales = hmask + nb * (QK_K / 8);
+ const sycl::half * d = (const sycl::half *) (scales + nb * 12);
+ return { qs, hmask, scales, d };
+ }
+
+ // unpack the 12 packed bytes into 16 6-bit scale codes (dequantize_row_q3_K
+ // aux layout), returned as float scale = d * (code - 32).
+ // done with wide (8/16-lane) ops rather than four 4-lane groups.
+ static ESIMD_INLINE sycl::ext::intel::esimd::simd<float, 16> unpack_scales(
+ sycl::ext::intel::esimd::simd<uint8_t, 12> in, float d) {
+ using namespace sycl::ext::intel::esimd;
+
+ // low 6-bit part: codes 0..7 = low nibble of bytes 0..7,
+ // codes 8..15 = high nibble of bytes 0..7
+ simd<uint8_t, 8> lo8 = in.select<8, 1>(0);
+ simd<uint8_t, 16> code;
+ code.select<8, 1>(0) = lo8 & simd<uint8_t, 8>(0x0F);
+ code.select<8, 1>(8) = lo8 >> simd<uint8_t, 8>(4);
+
+ // high 2-bit part: bytes 8..11 replicated 4x, group g (0..3) shifted 2*g
+ simd<uint8_t, 16> hib;
+ hib.select<4, 1>(0) = in.select<4, 1>(8);
+ hib.select<4, 1>(4) = in.select<4, 1>(8);
+ hib.select<4, 1>(8) = in.select<4, 1>(8);
+ hib.select<4, 1>(12) = in.select<4, 1>(8);
+ simd<uint8_t, 16> hshift;
+ hshift.select<4, 1>(0) = 0;
+ hshift.select<4, 1>(4) = 2;
+ hshift.select<4, 1>(8) = 4;
+ hshift.select<4, 1>(12) = 6;
+ hib = (hib >> hshift) & simd<uint8_t, 16>(0x03);
+
+ code = code | (hib << simd<uint8_t, 16>(4));
+ return (convert<float>(code) - 32.0f) * d;
+ }
+
+ 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<float, 256> & y_vec,
+ sycl::ext::intel::esimd::simd<float, 32> & acc_a,
+ sycl::ext::intel::esimd::simd<float, 32> & acc_b) {
+ using namespace sycl::ext::intel::esimd;
+
+ simd<uint8_t, 64> qs_a = block_load<uint8_t, 64>(pa.qs + bia * (QK_K / 4));
+ simd<uint8_t, 64> qs_b = 0;
+ simd<uint8_t, 32> hmask_a = block_load<uint8_t, 32>(pa.hmask + bia * (QK_K / 8));
+ simd<uint8_t, 32> hmask_b = 0;
+ simd<uint8_t, 12> scales_a = block_load<uint8_t, 12>(pa.scales + bia * 12);
+ simd<uint8_t, 12> scales_b = 0;
+
+ const float d_a = (float) pa.d[bia];
+ float d_b = 0.0f;
+ if (has_b) {
+ qs_b = block_load<uint8_t, 64>(pb.qs + bib * (QK_K / 4));
+ hmask_b = block_load<uint8_t, 32>(pb.hmask + bib * (QK_K / 8));
+ scales_b = block_load<uint8_t, 12>(pb.scales + bib * 12);
+ d_b = (float) pb.d[bib];
+ }
+
+ simd<float, 16> scale_f_a = unpack_scales(scales_a, d_a);
+ simd<float, 16> scale_f_b = unpack_scales(scales_b, d_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<float, 32> y_s = y_vec.select<32, 1>(s * 32);
+
+ // 2 low bits from qs, high bit from hmask (bit s of the same 32 bytes);
+ // value = (q & 3) + 4*bit - 4 (inverted hmask: subtract 4 when bit clear).
+ // merge in the integer domain: q3 = (q & 3) | (bit << 2) in {0..7},
+ // then a single convert + subtract yields q3 - 4 (one convert, not two)
+ simd<uint16_t, 32> q3_a = convert<uint16_t>(
+ (qs_a.select<32, 1>(byte_base) >> shift) & simd<uint8_t, 32>(3));
+ q3_a |= convert<uint16_t>(
+ ((hmask_a >> simd<uint8_t, 32>((uint8_t) s)) & simd<uint8_t, 32>(1)) << simd<uint8_t, 32>(2));
+ simd<uint16_t, 32> q3_b = convert<uint16_t>(
+ (qs_b.select<32, 1>(byte_base) >> shift) & simd<uint8_t, 32>(3));
+ q3_b |= convert<uint16_t>(
+ ((hmask_b >> simd<uint8_t, 32>((uint8_t) s)) & simd<uint8_t, 32>(1)) << simd<uint8_t, 32>(2));
+
+ simd<float, 32> qf_a = convert<float>(q3_a) - 4.0f;
+ simd<float, 32> qf_b = convert<float>(q3_b) - 4.0f;
+
+ const float scale_a_lo = scale_f_a[2 * s + 0];
+ const float scale_a_hi = scale_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];
+
+ simd<float, 32> scale_vec_a = splat_lo_hi(scale_a_lo, scale_a_hi);
+ simd<float, 32> scale_vec_b = splat_lo_hi(scale_b_lo, scale_b_hi);
+
+ simd<float, 32> deq_a = qf_a * scale_vec_a;
+ simd<float, 32> deq_b = qf_b * scale_vec_b;
+
+ acc_a += y_s * deq_a;
+ acc_b += y_s * deq_b;
+ }
+ }
+};
+
+// ---------------------------------------------------------------------------
+// Q4_K, SOA reorder layout produced by reorder_qw_q4_k:
+// [qs: nb*(QK_K/2)] [scales: nb*K_SCALE_SIZE] [dm: nb*sizeof(half2)]
+// with nb = nrows*num_blocks_per_row.
+// ---------------------------------------------------------------------------
+template <> struct esimd_reorder_q_traits<GGML_TYPE_Q4_K> {
+ 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 / 2);
+ const sycl::half * dm = (const sycl::half *) (scales + nb * K_SCALE_SIZE);
+ 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<float, 256> & y_vec,
+ sycl::ext::intel::esimd::simd<float, 32> & acc_a,
+ sycl::ext::intel::esimd::simd<float, 32> & acc_b) {
+ using namespace sycl::ext::intel::esimd;
+
+ simd<uint8_t, 128> qs_a = block_load<uint8_t, 128>(pa.qs + bia * (QK_K / 2));
+ simd<uint8_t, 128> qs_b = 0;
+ simd<uint8_t, 12> scales_a = block_load<uint8_t, 12>(pa.scales + bia * K_SCALE_SIZE);
+ simd<uint8_t, 12> 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<uint8_t, 128>(pb.qs + bib * (QK_K / 2));
+ scales_b = block_load<uint8_t, 12>(pb.scales + bib * K_SCALE_SIZE);
+ dall_b = (float) pb.dm[bib * 2 + 0];
+ dmin_b = (float) pb.dm[bib * 2 + 1];
+ }
+
+ simd<float, 8> scale_f_a, min_f_a, scale_f_b, min_f_b;
+ unpack_scale_min_k4(scales_a, dall_a, dmin_a, scale_f_a, min_f_a);
+ unpack_scale_min_k4(scales_b, dall_b, dmin_b, scale_f_b, min_f_b);
+
+ simd<uint8_t, 128> qs_lo_a = qs_a & simd<uint8_t, 128>(0x0F);
+ simd<uint8_t, 128> qs_hi_a = qs_a >> simd<uint8_t, 128>(4);
+ simd<uint8_t, 128> qs_lo_b = qs_b & simd<uint8_t, 128>(0x0F);
+ simd<uint8_t, 128> qs_hi_b = qs_b >> simd<uint8_t, 128>(4);
+
+#pragma unroll
+ for (int sb = 0; sb < 8; sb += 2) {
+ const int q_offset = sb * 16;
+ simd<float, 32> y_lo = y_vec.select<32, 1>(sb * 32);
+ simd<float, 32> y_hi = y_vec.select<32, 1>((sb + 1) * 32);
+
+ const float scale_a_lo = scale_f_a[sb];
+ const float scale_a_hi = scale_f_a[sb + 1];
+ const float min_a_lo = min_f_a[sb];
+ const float min_a_hi = min_f_a[sb + 1];
+ const float scale_b_lo = scale_f_b[sb];
+ const float scale_b_hi = scale_f_b[sb + 1];
+ const float min_b_lo = min_f_b[sb];
+ const float min_b_hi = min_f_b[sb + 1];
+
+ simd<uint8_t, 32> qa_lo = qs_lo_a.select<32, 1>(q_offset);
+ simd<uint8_t, 32> qa_hi = qs_hi_a.select<32, 1>(q_offset);
+ simd<uint8_t, 32> qb_lo = qs_lo_b.select<32, 1>(q_offset);
+ simd<uint8_t, 32> qb_hi = qs_hi_b.select<32, 1>(q_offset);
+
+ simd<float, 32> deq_a_lo = convert<float>(qa_lo) * scale_a_lo + min_a_lo;
+ simd<float, 32> deq_a_hi = convert<float>(qa_hi) * scale_a_hi + min_a_hi;
+ simd<float, 32> deq_b_lo = convert<float>(qb_lo) * scale_b_lo + min_b_lo;
+ simd<float, 32> deq_b_hi = convert<float>(qb_hi) * scale_b_hi + min_b_hi;
+
+ acc_a += y_lo * deq_a_lo;
+ acc_b += y_lo * deq_b_lo;
+ acc_a += y_hi * deq_a_hi;
+ acc_b += y_hi * deq_b_hi;
+ }
+ }
+};
+
+// ---------------------------------------------------------------------------
+// Q6_K, SOA reorder layout:
+// [ql: nb*(QK_K/2)] [qh: nb*(QK_K/4)] [scales(int8): nb*(QK_K/16)] [d: nb*half]
+// ---------------------------------------------------------------------------
+template <> struct esimd_reorder_q_traits<GGML_TYPE_Q6_K> {
+ struct ptrs {
+ const uint8_t * ql;
+ const uint8_t * qh;
+ const int8_t * scales;
+ const sycl::half * d;
+ };
+
+ static ESIMD_INLINE ptrs make_ptrs(const void * vx, size_t nb) {
+ const uint8_t * ql = (const uint8_t *) vx;
+ const uint8_t * qh = ql + nb * (QK_K / 2);
+ const int8_t * scales = (const int8_t *) (qh + nb * (QK_K / 4));
+ const sycl::half * d = (const sycl::half *) (scales + nb * (QK_K / 16));
+ return { ql, qh, scales, d };
+ }
+
+ 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<float, 256> & y_vec,
+ sycl::ext::intel::esimd::simd<float, 32> & acc_a,
+ sycl::ext::intel::esimd::simd<float, 32> & acc_b) {
+ using namespace sycl::ext::intel::esimd;
+
+ simd<uint8_t, 128> ql_a = block_load<uint8_t, 128>(pa.ql + bia * (QK_K / 2));
+ simd<uint8_t, 128> ql_b = 0;
+ simd<uint8_t, 64> qh_a = block_load<uint8_t, 64>(pa.qh + bia * (QK_K / 4));
+ simd<uint8_t, 64> qh_b = 0;
+ simd<int8_t, 16> scales_a = block_load<int8_t, 16>(pa.scales + bia * (QK_K / 16));
+ simd<int8_t, 16> scales_b = 0;
+
+ const float d_a = (float) pa.d[bia];
+ float d_b = 0.0f;
+ if (has_b) {
+ ql_b = block_load<uint8_t, 128>(pb.ql + bib * (QK_K / 2));
+ qh_b = block_load<uint8_t, 64>(pb.qh + bib * (QK_K / 4));
+ scales_b = block_load<int8_t, 16>(pb.scales + bib * (QK_K / 16));
+ d_b = (float) pb.d[bib];
+ }
+
+ simd<float, 16> sc_a = convert<float>(scales_a);
+ simd<float, 16> sc_b = convert<float>(scales_b);
+
+#pragma unroll
+ for (int im = 0; im < 2; ++im) {
+ simd<uint8_t, 32> ql_lo_a = ql_a.select<32, 1>(64 * im);
+ simd<uint8_t, 32> ql_hi_a = ql_a.select<32, 1>(64 * im + 32);
+ simd<uint8_t, 32> qh_bits_a = qh_a.select<32, 1>(32 * im);
+ simd<uint8_t, 32> ql_lo_b = ql_b.select<32, 1>(64 * im);
+ simd<uint8_t, 32> ql_hi_b = ql_b.select<32, 1>(64 * im + 32);
+ simd<uint8_t, 32> qh_bits_b = qh_b.select<32, 1>(32 * im);
+
+ // reconstruct each 32-wide 6-bit group (matches dequantize_row_q6_K)
+#pragma unroll
+ for (int g = 0; g < 4; ++g) {
+ simd<float, 32> y_g = y_vec.select<32, 1>(32 * (4 * im + g));
+
+ const float scale_a_lo = sc_a[8 * im + 2 * g + 0] * d_a;
+ const float scale_a_hi = sc_a[8 * im + 2 * g + 1] * d_a;
+ const float scale_b_lo = sc_b[8 * im + 2 * g + 0] * d_b;
+ const float scale_b_hi = sc_b[8 * im + 2 * g + 1] * d_b;
+
+ simd<float, 32> scale_vec_a = splat_lo_hi(scale_a_lo, scale_a_hi);
+ simd<float, 32> scale_vec_b = splat_lo_hi(scale_b_lo, scale_b_hi);
+
+ simd<uint8_t, 32> qa;
+ simd<uint8_t, 32> qb;
+ switch (g) {
+ case 0:
+ qa = (ql_lo_a & simd<uint8_t, 32>(0x0F)) | ((qh_bits_a & simd<uint8_t, 32>(0x03)) << simd<uint8_t, 32>(4));
+ qb = (ql_lo_b & simd<uint8_t, 32>(0x0F)) | ((qh_bits_b & simd<uint8_t, 32>(0x03)) << simd<uint8_t, 32>(4));
+ break;
+ case 1:
+ qa = (ql_hi_a & simd<uint8_t, 32>(0x0F)) | ((qh_bits_a & simd<uint8_t, 32>(0x0C)) << simd<uint8_t, 32>(2));
+ qb = (ql_hi_b & simd<uint8_t, 32>(0x0F)) | ((qh_bits_b & simd<uint8_t, 32>(0x0C)) << simd<uint8_t, 32>(2));
+ break;
+ case 2:
+ qa = (ql_lo_a >> simd<uint8_t, 32>(4)) | (qh_bits_a & simd<uint8_t, 32>(0x30));
+ qb = (ql_lo_b >> simd<uint8_t, 32>(4)) | (qh_bits_b & simd<uint8_t, 32>(0x30));
+ break;
+ default:
+ qa = (ql_hi_a >> simd<uint8_t, 32>(4)) | ((qh_bits_a & simd<uint8_t, 32>(0xC0)) >> simd<uint8_t, 32>(2));
+ qb = (ql_hi_b >> simd<uint8_t, 32>(4)) | ((qh_bits_b & simd<uint8_t, 32>(0xC0)) >> simd<uint8_t, 32>(2));
+ break;
+ }
+
+ simd<float, 32> deq_a = (convert<float>(qa) - 32.0f) * scale_vec_a;
+ simd<float, 32> deq_b = (convert<float>(qb) - 32.0f) * scale_vec_b;
+
+ acc_a += y_g * deq_a;
+ acc_b += y_g * deq_b;
+ }
+ }
+ }
+};
+
+} // namespace ggml_sycl_esimd
+
+#endif // GGML_SYCL_ESIMD_HPP
overview / pkg / ggml / sources / llama.cpp / commitdiff