Add Q5_K ESIMD kernel

ggml/src/ggml-sycl/dmmv.cpp

@@ -1955,6 +1955,23 @@ static void dequantize_mul_mat_vec_q4_K_sycl_reorder_esimd(const void *vx, const
     });
 }
 
+static void dequantize_mul_mat_vec_q5_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_Q5_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,
@@ -2134,7 +2151,15 @@ void ggml_sycl_op_dequantize_mul_mat_vec(
         case GGML_TYPE_Q5_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_q5_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_q5_K_sycl_reorder_esimd(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+                }
+                else
+#endif
+                {
+                    dequantize_mul_mat_vec_q5_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
+                }
             } else {
                 dequantize_mul_mat_vec_q5_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream);
             }

ggml/src/ggml-sycl/esimd.hpp

@@ -1,15 +1,3 @@
-//
-// 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
 
@@ -287,6 +275,128 @@ template <> struct esimd_reorder_q_traits<GGML_TYPE_Q4_K> {
     }
 };
 
+// ---------------------------------------------------------------------------
+// Q5_K, SOA reorder layout produced by reorder_qw_q5_k:
+//   [qs: nb*(QK_K/2)] [qh: nb*(QK_K/8)] [scales: nb*K_SCALE_SIZE] [dm: nb*sizeof(half2)]
+// with nb = nrows*num_blocks_per_row.
+//
+// Identical to Q4_K except each 4-bit quant gains a 5th (high) bit from qh:
+// output chunk c (0..7) adds 16 when bit c of qh[l] is set, where qh[l] indexes
+// the same 32 bytes for every chunk (matches dequantize_row_q5_K).
+// ---------------------------------------------------------------------------
+template <> struct esimd_reorder_q_traits<GGML_TYPE_Q5_K> {
+    struct ptrs {
+        const uint8_t *    qs;
+        const uint8_t *    qh;
+        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 * qh     = qs + nb * (QK_K / 2);
+        const uint8_t * scales = qh + nb * (QK_K / 8);
+        const sycl::half * dm  = (const sycl::half *) (scales + nb * K_SCALE_SIZE);
+        return { qs, qh, scales, dm };
+    }
+
+    // extract bit `bit` (0..7) of each lane and move it to bit position 4,
+    // e.g. for the 4-bit base quant's 5th (high) bit. `bit` is always a
+    // compile-time-known unrolled loop constant at call sites, so this folds
+    // to a single mask (bit==4), mask+left-shift (bit<4), or mask+right-shift
+    // (bit>4) instead of the shift+mask+shift a naive `(qh>>bit & 1) << 4` emits.
+    static ESIMD_INLINE sycl::ext::intel::esimd::simd<uint16_t, 32> extract_bit_to_pos4(
+            sycl::ext::intel::esimd::simd<uint8_t, 32> qh, int bit) {
+        using namespace sycl::ext::intel::esimd;
+        simd<uint16_t, 32> masked = convert<uint16_t>(qh & simd<uint8_t, 32>((uint8_t) (1u << bit)));
+        if (bit < 4) {
+            return masked << simd<uint16_t, 32>((uint16_t) (4 - bit));
+        } else if (bit > 4) {
+            return masked >> simd<uint16_t, 32>((uint16_t) (bit - 4));
+        }
+        return masked;
+    }
+
+    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, 32>  qh_a     = block_load<uint8_t, 32>(pa.qh + bia * (QK_K / 8));
+        simd<uint8_t, 32>  qh_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));
+            qh_b     = block_load<uint8_t, 32>(pb.qh + bib * (QK_K / 8));
+            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_u8 = qs_lo_a.select<32, 1>(q_offset);
+            simd<uint8_t, 32> qa_hi_u8 = qs_hi_a.select<32, 1>(q_offset);
+            simd<uint8_t, 32> qb_lo_u8 = qs_lo_b.select<32, 1>(q_offset);
+            simd<uint8_t, 32> qb_hi_u8 = qs_hi_b.select<32, 1>(q_offset);
+            simd<uint16_t, 32> qa_lo = convert<uint16_t>(qa_lo_u8);
+            simd<uint16_t, 32> qa_hi = convert<uint16_t>(qa_hi_u8);
+            simd<uint16_t, 32> qb_lo = convert<uint16_t>(qb_lo_u8);
+            simd<uint16_t, 32> qb_hi = convert<uint16_t>(qb_hi_u8);
+
+            // add the 5th bit: chunk sb uses qh bit sb, chunk sb+1 uses qh bit sb+1;
+            // qh always indexes the same 32 bytes regardless of chunk
+            qa_lo += extract_bit_to_pos4(qh_a, sb);
+            qa_hi += extract_bit_to_pos4(qh_a, sb + 1);
+            qb_lo += extract_bit_to_pos4(qh_b, sb);
+            qb_hi += extract_bit_to_pos4(qh_b, sb + 1);
+
+            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]

ggml/src/ggml-sycl/ggml-sycl.cpp

@@ -3750,6 +3750,7 @@ static bool ggml_sycl_supports_reorder_esimd(enum ggml_type type) {
     switch (type) {
         case GGML_TYPE_Q3_K:
         case GGML_TYPE_Q4_K:
+        case GGML_TYPE_Q5_K:
         case GGML_TYPE_Q6_K:
             return true;
         default: