diff --git a/docs/backend/SYCL.md b/docs/backend/SYCL.md
index 73f89a706..519d03975 100644
--- a/docs/backend/SYCL.md
+++ b/docs/backend/SYCL.md
@@ -804,6 +804,7 @@ User can use the device management in [docs/multi-gpu.md](https://github.com/ggm
| GGML_SYCL_MKL_FA_DEBUG | 0 (default) or 1 | Enable per-call diagnostic logging for MKL flash attention: GEMM/softmax timings, interleaved-head detection, and buffer memory usage. |
| GGML_SYCL_MKL_FA_DIAG | 0 (default) or 1 | Enable output fingerprinting for MKL flash attention. Dumps the first 64 float output values for the first 6 FA calls with n_kv ≥ 1024, labeled with kernel type (MKL/TILE/VEC) for cross-kernel comparison. |
| GGML_SYCL_ENABLE_FUSION | 0 or 1 (default) | Enable fused-kernel dispatch in graph compute (currently top-k MoE gating). |
+| GGML_SYCL_ENABLE_ESIMD | 0 or 1 (default)| Enable ESIMD kernels when available. |
| ZES_ENABLE_SYSMAN | 0 (default) or 1 | Support to get free memory of GPU by sycl::aspect::ext_intel_free_memory.
Recommended to use when --split-mode = layer |
| UR_L0_ENABLE_RELAXED_ALLOCATION_LIMITS | 0 (default) or 1 | Allow SYCL/Unified Runtime Level Zero device allocations larger than 4 GiB. llama.cpp's direct Level Zero allocation path requests the relaxed maximum-size limit itself when GGML_SYCL_ENABLE_LEVEL_ZERO=1. |
| GGML_SYCL_USM_SYSTEM | 0 (default) or 1 | Enable experimental support for [USM system allocations](https://github.khronos.org/SYCL_Reference/iface/usm_basic_concept.html#system-allocations) for large GPU buffers. This requires enough host memory for model weights and caches, an Intel Xe2+ GPU such as BMG or newer and supported on Linux only, with CONFIG_DRM_XE_GPUSVM enabled. |
diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp
index 4fa34a526..34de284d8 100644
--- a/ggml/src/ggml-sycl/common.hpp
+++ b/ggml/src/ggml-sycl/common.hpp
@@ -61,6 +61,7 @@ void ggml_sycl_host_free(void* ptr);
extern int g_ggml_sycl_debug;
extern int g_ggml_sycl_enable_optimize;
extern int g_ggml_sycl_enable_fusion;
+extern int g_ggml_sycl_enable_esimd;
extern int g_ggml_sycl_prioritize_dmmv;
extern int g_ggml_sycl_enable_flash_attention;
extern int g_ggml_sycl_dev2dev_memcpy;
diff --git a/ggml/src/ggml-sycl/dmmv.cpp b/ggml/src/ggml-sycl/dmmv.cpp
index ee7cd2d48..d8da0a16b 100644
--- a/ggml/src/ggml-sycl/dmmv.cpp
+++ b/ggml/src/ggml-sycl/dmmv.cpp
@@ -8,6 +8,9 @@
#include
#define GGML_SYCL_DMMV_HAS_BF16
#endif
+ #include
+ #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){
@@ -1864,6 +1867,113 @@ static void dequantize_mul_mat_vec_q6_K_sycl(const void *vx, const float *y,
});
}
+#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
+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 lmem,
+ const sycl::nd_item<1> & it) {
+ using namespace sycl::ext::intel::esimd;
+ using traits = ggml_sycl_esimd::esimd_reorder_q_traits;
+
+ 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 acc0 = 0.0f;
+ simd acc1 = 0.0f;
+
+ for (int ib = tid; ib < num_blocks_per_row; ib += GGML_SYCL_DMMV_ESIMD_WG_SIZE) {
+ simd y_vec = block_load(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(acc0, std::plus<>{});
+ lmem[tid * 2 + 1] = reduce(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 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_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 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_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 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);
+ });
+ });
+}
+
+#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,
@@ -1992,7 +2102,15 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
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);
}
@@ -2000,7 +2118,15 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
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);
}
@@ -2016,7 +2142,15 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
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);
}
diff --git a/ggml/src/ggml-sycl/esimd.hpp b/ggml/src/ggml-sycl/esimd.hpp
new file mode 100644
index 000000000..d7609b11f
--- /dev/null
+++ b/ggml/src/ggml-sycl/esimd.hpp
@@ -0,0 +1,392 @@
+//
+// 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
+
+#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::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 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 splat_lo_hi(float lo, float hi) {
+ using namespace sycl::ext::intel::esimd;
+ simd 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 scales, float dall, float dmin,
+ sycl::ext::intel::esimd::simd & scale_f,
+ sycl::ext::intel::esimd::simd & min_f) {
+ using namespace sycl::ext::intel::esimd;
+ simd sc = 0;
+ simd m = 0;
+ simd scale_lo = scales.select<4, 1>(0);
+ simd min_lo = scales.select<4, 1>(4);
+ simd hi_bits = scales.select<4, 1>(8);
+ sc.select<4, 1>(0) = scale_lo & simd(0x3F);
+ sc.select<4, 1>(4) = (hi_bits & simd(0x0F)) |
+ ((scale_lo >> simd(6)) << simd(4));
+ m.select<4, 1>(0) = min_lo & simd(0x3F);
+ m.select<4, 1>(4) = (hi_bits >> simd(4)) |
+ ((min_lo >> simd(6)) << simd(4));
+ scale_f = convert(sc) * dall;
+ min_f = convert(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 {
+ 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 unpack_scales(
+ sycl::ext::intel::esimd::simd 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 lo8 = in.select<8, 1>(0);
+ simd code;
+ code.select<8, 1>(0) = lo8 & simd(0x0F);
+ code.select<8, 1>(8) = lo8 >> simd(4);
+
+ // high 2-bit part: bytes 8..11 replicated 4x, group g (0..3) shifted 2*g
+ simd 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 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(0x03);
+
+ code = code | (hib << simd(4));
+ return (convert(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 & 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 hmask_a = block_load(pa.hmask + bia * (QK_K / 8));
+ simd hmask_b = 0;
+ simd scales_a = block_load(pa.scales + bia * 12);
+ simd scales_b = 0;
+
+ const float d_a = (float) pa.d[bia];
+ float d_b = 0.0f;
+ if (has_b) {
+ qs_b = block_load(pb.qs + bib * (QK_K / 4));
+ hmask_b = block_load(pb.hmask + bib * (QK_K / 8));
+ scales_b = block_load(pb.scales + bib * 12);
+ d_b = (float) pb.d[bib];
+ }
+
+ simd scale_f_a = unpack_scales(scales_a, d_a);
+ simd 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 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 q3_a = convert(
+ (qs_a.select<32, 1>(byte_base) >> shift) & simd(3));
+ q3_a |= convert(
+ ((hmask_a >> simd((uint8_t) s)) & simd(1)) << simd(2));
+ simd q3_b = convert(
+ (qs_b.select<32, 1>(byte_base) >> shift) & simd(3));
+ q3_b |= convert(
+ ((hmask_b >> simd((uint8_t) s)) & simd(1)) << simd(2));
+
+ simd qf_a = convert(q3_a) - 4.0f;
+ simd qf_b = convert(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 scale_vec_a = splat_lo_hi(scale_a_lo, scale_a_hi);
+ simd scale_vec_b = splat_lo_hi(scale_b_lo, scale_b_hi);
+
+ simd deq_a = qf_a * scale_vec_a;
+ simd 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 {
+ 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 & 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 / 2));
+ simd qs_b = 0;
+ simd scales_a = block_load(pa.scales + bia * K_SCALE_SIZE);
+ 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 / 2));
+ scales_b = block_load(pb.scales + bib * K_SCALE_SIZE);
+ dall_b = (float) pb.dm[bib * 2 + 0];
+ dmin_b = (float) pb.dm[bib * 2 + 1];
+ }
+
+ simd 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 qs_lo_a = qs_a & simd(0x0F);
+ simd qs_hi_a = qs_a >> simd(4);
+ simd qs_lo_b = qs_b & simd(0x0F);
+ simd qs_hi_b = qs_b >> simd(4);
+
+#pragma unroll
+ for (int sb = 0; sb < 8; sb += 2) {
+ const int q_offset = sb * 16;
+ simd y_lo = y_vec.select<32, 1>(sb * 32);
+ simd 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 qa_lo = qs_lo_a.select<32, 1>(q_offset);
+ simd qa_hi = qs_hi_a.select<32, 1>(q_offset);
+ simd qb_lo = qs_lo_b.select<32, 1>(q_offset);
+ simd qb_hi = qs_hi_b.select<32, 1>(q_offset);
+
+ simd deq_a_lo = convert(qa_lo) * scale_a_lo + min_a_lo;
+ simd deq_a_hi = convert(qa_hi) * scale_a_hi + min_a_hi;
+ simd deq_b_lo = convert(qb_lo) * scale_b_lo + min_b_lo;
+ simd deq_b_hi = convert(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 {
+ 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 & y_vec,
+ sycl::ext::intel::esimd::simd & acc_a,
+ sycl::ext::intel::esimd::simd & acc_b) {
+ using namespace sycl::ext::intel::esimd;
+
+ simd ql_a = block_load(pa.ql + bia * (QK_K / 2));
+ simd ql_b = 0;
+ simd qh_a = block_load(pa.qh + bia * (QK_K / 4));
+ simd qh_b = 0;
+ simd scales_a = block_load(pa.scales + bia * (QK_K / 16));
+ simd scales_b = 0;
+
+ const float d_a = (float) pa.d[bia];
+ float d_b = 0.0f;
+ if (has_b) {
+ ql_b = block_load(pb.ql + bib * (QK_K / 2));
+ qh_b = block_load(pb.qh + bib * (QK_K / 4));
+ scales_b = block_load(pb.scales + bib * (QK_K / 16));
+ d_b = (float) pb.d[bib];
+ }
+
+ simd sc_a = convert(scales_a);
+ simd sc_b = convert(scales_b);
+
+#pragma unroll
+ for (int im = 0; im < 2; ++im) {
+ simd ql_lo_a = ql_a.select<32, 1>(64 * im);
+ simd ql_hi_a = ql_a.select<32, 1>(64 * im + 32);
+ simd qh_bits_a = qh_a.select<32, 1>(32 * im);
+ simd ql_lo_b = ql_b.select<32, 1>(64 * im);
+ simd ql_hi_b = ql_b.select<32, 1>(64 * im + 32);
+ simd 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 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 scale_vec_a = splat_lo_hi(scale_a_lo, scale_a_hi);
+ simd scale_vec_b = splat_lo_hi(scale_b_lo, scale_b_hi);
+
+ simd qa;
+ simd qb;
+ switch (g) {
+ case 0:
+ qa = (ql_lo_a & simd(0x0F)) | ((qh_bits_a & simd(0x03)) << simd(4));
+ qb = (ql_lo_b & simd(0x0F)) | ((qh_bits_b & simd(0x03)) << simd(4));
+ break;
+ case 1:
+ qa = (ql_hi_a & simd(0x0F)) | ((qh_bits_a & simd(0x0C)) << simd(2));
+ qb = (ql_hi_b & simd(0x0F)) | ((qh_bits_b & simd(0x0C)) << simd(2));
+ break;
+ case 2:
+ qa = (ql_lo_a >> simd(4)) | (qh_bits_a & simd(0x30));
+ qb = (ql_lo_b >> simd(4)) | (qh_bits_b & simd(0x30));
+ break;
+ default:
+ qa = (ql_hi_a >> simd(4)) | ((qh_bits_a & simd(0xC0)) >> simd(2));
+ qb = (ql_hi_b >> simd(4)) | ((qh_bits_b & simd(0xC0)) >> simd(2));
+ break;
+ }
+
+ simd deq_a = (convert(qa) - 32.0f) * scale_vec_a;
+ simd deq_b = (convert(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
diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp
index 8a19f648b..b98a28832 100644
--- a/ggml/src/ggml-sycl/ggml-sycl.cpp
+++ b/ggml/src/ggml-sycl/ggml-sycl.cpp
@@ -43,6 +43,9 @@
# include
# define GGML_SYCL_SUPPORT_VMM
#endif
+#if defined(__INTEL_LLVM_COMPILER)
+ #define GGML_SYCL_DMMV_HAS_ESIMD
+#endif
#include
#include "ggml.h"
@@ -90,6 +93,7 @@ int g_ggml_sycl_fa_onednn = 1;
int g_ggml_sycl_fa_onednn_max_kv = 0;
int g_ggml_sycl_enable_vmm = 1;
int g_ggml_sycl_enable_fusion = 1;
+int g_ggml_sycl_enable_esimd = 1;
int g_ggml_sycl_prioritize_dmmv = 0;
int g_ggml_sycl_use_async_mem_op = 0;
int g_ggml_sycl_use_async_mem_op_requested = 1;
@@ -298,6 +302,7 @@ static void ggml_check_sycl() try {
g_ggml_sycl_fa_onednn_max_kv = ggml_sycl_get_env("GGML_SYCL_FA_ONEDNN_MAX_KV", 0);
g_ggml_sycl_enable_vmm = ggml_sycl_get_env("GGML_SYCL_ENABLE_VMM", 1);
g_ggml_sycl_enable_fusion = ggml_sycl_get_env("GGML_SYCL_ENABLE_FUSION", 1);
+ g_ggml_sycl_enable_esimd = ggml_sycl_get_env("GGML_SYCL_ENABLE_ESIMD", 1);
g_ggml_sycl_prioritize_dmmv = ggml_sycl_get_env("GGML_SYCL_PRIORITIZE_DMMV", 0);
g_ggml_sycl_dev2dev_memcpy = ggml_sycl_get_env("GGML_SYCL_DEV2DEV_MEMCPY", DEV2DEV_MEMCPY_SYCL);
@@ -392,6 +397,12 @@ static void ggml_check_sycl() try {
GGML_LOG_INFO(" GGML_SYCL_ENABLE_FUSION: %d\n", g_ggml_sycl_enable_fusion);
+#if defined(__INTEL_LLVM_COMPILER)
+ GGML_LOG_INFO(" GGML_SYCL_ENABLE_ESIMD: %d\n", g_ggml_sycl_enable_esimd);
+#else
+ GGML_LOG_INFO(" GGML_SYCL_ENABLE_ESIMD: %d disabled by compile flag\n", g_ggml_sycl_enable_esimd);
+#endif
+
GGML_LOG_INFO(" GGML_SYCL_PRIORITIZE_DMMV: %d\n", g_ggml_sycl_prioritize_dmmv);
g_ggml_sycl_use_async_mem_op_requested = ggml_sycl_get_env("GGML_SYCL_USE_ASYNC_MEM_OP", 1);
@@ -3740,6 +3751,22 @@ 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_Q3_K:
+ case GGML_TYPE_Q4_K:
+ case GGML_TYPE_Q6_K:
+ return true;
+ default:
+ return false;
+ }
+#else
+ GGML_UNUSED(type);
+ return false;
+#endif
+}
+
static bool ggml_sycl_supports_dmmv(enum ggml_type type) {
switch (type) {
case GGML_TYPE_Q1_0:
@@ -4443,19 +4470,22 @@ static void ggml_sycl_mul_mat(ggml_backend_sycl_context & ctx, const ggml_tensor
use_mul_mat_q = use_mul_mat_q && (src1->ne[1] <= MMQ_MAX_BATCH_SIZE);
#endif // SYCL_USE_XMX
- // Dispatch becomes obscure with the reorder, MMVQ when the reorder optimization
- // is enabled takes precedence over DMMV, the current if-else implementation
- // requires disabling DMMV if both conditions are met
+ // When reorder is enabled, both ESIMD, MMVQ and DMMV kernels may be used. For
+ // best performance use ESIMD when supported, followed by MMVQ, and finally DMMV.
+ // But the reordered ESIMD path cannot be used without reordered MMVQ. A later
+ // multi-token call (ne[1] in 2..8) will take the MMVQ path and it would read the
+ // reordered bytes as if they were still the unreordered layout.
if (!g_ggml_sycl_prioritize_dmmv && ((should_reorder_tensor(ctx, dst) &&
ggml_sycl_supports_reorder_mmvq(src0->type)))) {
- // Arc770 get benefit with Q4_0 by skipping it.
- if (!(ggml_sycl_info().devices[ctx.device].hw_info.arch ==
- gpu_arch::intel_gpu_acm_g10 &&
- src0->type == GGML_TYPE_Q4_0)) {
- use_dequantize_mul_mat_vec =
- use_dequantize_mul_mat_vec && !use_mul_mat_vec_q;
- }
+ bool use = g_ggml_sycl_enable_esimd && ggml_sycl_supports_reorder_esimd(src0->type);
+ // Arc770 get benefit with Q4_0 by skipping MMVQ path
+ if (!(ggml_sycl_info().devices[ctx.device].hw_info.arch ==
+ gpu_arch::intel_gpu_acm_g10 &&
+ src0->type == GGML_TYPE_Q4_0)) {
+ use = use || !use_mul_mat_vec_q;
+ }
+ use_dequantize_mul_mat_vec = use_dequantize_mul_mat_vec && use;
}
if (!split && src0->type == GGML_TYPE_F16 && ggml_is_permuted(src0) && ggml_is_permuted(src1) && src1->ne[1] == 1) {