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) {