From de73a46a9a3cb22b37b9ad2d25bb51a918e42986 Mon Sep 17 00:00:00 2001 From: Todd Malsbary Date: Wed, 15 Jul 2026 21:44:06 +0000 Subject: [PATCH 1/8] Add DMMV Q4_K and Q6_K ESIMD kernels Configure cmake build with -DGGML_SYCL_ESIMD=ON to enable. Signed-off-by: Todd Malsbary --- docs/backend/SYCL.md | 1 + ggml/CMakeLists.txt | 1 + ggml/src/ggml-sycl/CMakeLists.txt | 4 + ggml/src/ggml-sycl/dmmv.cpp | 309 ++++++++++++++++++++++++++++++ 4 files changed, 315 insertions(+) diff --git a/docs/backend/SYCL.md b/docs/backend/SYCL.md index 0814ceb60f92..934f5bef01db 100644 --- a/docs/backend/SYCL.md +++ b/docs/backend/SYCL.md @@ -776,6 +776,7 @@ use 1 SYCL GPUs: [0] with Max compute units:512 | GGML_SYCL_F16 | OFF *(default)* \|ON *(optional)* | Enable FP16 build with SYCL code path. (1.) | | GGML_SYCL_GRAPH | ON *(default)* \|OFF *(Optional)* | Enable build with [SYCL Graph extension](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc). | | GGML_SYCL_DNN | ON *(default)* \|OFF *(Optional)* | Enable build with oneDNN. | +| GGML_SYCL_ESIMD | OFF *(default)* \|ON *(Optional)* | Enable build with ESIMD kernels. | | GGML_SYCL_HOST_MEM_FALLBACK | ON *(default)* \|OFF *(Optional)* | Allow host memory fallback when device memory is full during quantized weight reorder. Enables inference to continue at reduced speed (reading over PCIe) instead of failing. Requires Linux kernel 6.8+. | | GGML_SYCL_SUPPORT_LEVEL_ZERO_API | ON *(default)* \|OFF *(Optional)* | Support to use Level Zero API for device memory allocation. Requires Level Zero headers/library at build time and Intel GPU driver (Level Zero runtime) at run time. Reduces system RAM usage during multi-GPU inference. SYCL backend always runs on Level Zero running time even if it's set as OFF (The SYCL api will be usage for memory allocation).| | CMAKE_C_COMPILER | `icx` *(Linux)*, `icx/cl` *(Windows)* | Set `icx` compiler for SYCL code path. | diff --git a/ggml/CMakeLists.txt b/ggml/CMakeLists.txt index 5381c2136203..ed059fbaafe0 100644 --- a/ggml/CMakeLists.txt +++ b/ggml/CMakeLists.txt @@ -251,6 +251,7 @@ option(GGML_SYCL_GRAPH "ggml: enable graphs in the SYCL bac option(GGML_SYCL_HOST_MEM_FALLBACK "ggml: allow host memory fallback in SYCL reorder (requires kernel 6.8+)" ON) option(GGML_SYCL_SUPPORT_LEVEL_ZERO_API "ggml: use Level Zero API in SYCL backend" ON) option(GGML_SYCL_DNN "ggml: enable oneDNN in the SYCL backend" ON) +option(GGML_SYCL_ESIMD "ggml: enable ESIMD kernels in the SYCL backend" OFF) set (GGML_SYCL_TARGET "INTEL" CACHE STRING "ggml: sycl target device") set (GGML_SYCL_DEVICE_ARCH "" CACHE STRING diff --git a/ggml/src/ggml-sycl/CMakeLists.txt b/ggml/src/ggml-sycl/CMakeLists.txt index 1c17d20df12b..332c57c0cf1a 100644 --- a/ggml/src/ggml-sycl/CMakeLists.txt +++ b/ggml/src/ggml-sycl/CMakeLists.txt @@ -158,6 +158,10 @@ else() endif() target_compile_definitions(ggml-sycl PRIVATE GGML_SYCL_DNNL=${GGML_SYCL_DNNL}) +if (GGML_SYCL_ESIMD) + add_compile_definitions(GGML_SYCL_ESIMD) +endif() + if (GGML_SYCL_F16) add_compile_definitions(GGML_SYCL_F16) endif() diff --git a/ggml/src/ggml-sycl/dmmv.cpp b/ggml/src/ggml-sycl/dmmv.cpp index fa62975d3ff5..529912801864 100644 --- a/ggml/src/ggml-sycl/dmmv.cpp +++ b/ggml/src/ggml-sycl/dmmv.cpp @@ -10,6 +10,13 @@ #endif #endif +#if defined(GGML_SYCL_ESIMD) + #if !defined(__INTEL_LLVM_COMPILER) + #error "GGML_SYCL_ESIMD requires the Intel oneAPI (DPC++) compiler" + #endif + #include +#endif + static void convert_f16(const void * vx, const int64_t ib, const int iqs, dfloat2 & v){ const sycl::half *x = (const sycl::half *)vx; @@ -1855,6 +1862,300 @@ static void dequantize_mul_mat_vec_q6_K_sycl(const void *vx, const float *y, }); } +#ifdef GGML_SYCL_ESIMD +constexpr int GGML_SYCL_DMMV_ESIMD_WG_SIZE = 4; + +ESIMD_INLINE static void dequantize_mul_mat_vec_q4_k_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; + + const int num_blocks_per_row = ncols / QK_K; + const size_t nb = (size_t)nrows * num_blocks_per_row; + const uint8_t * qs_base = (const uint8_t *)vx; + const uint8_t * scales_base = qs_base + nb * (QK_K / 2); + const sycl::half * dm_base = (const sycl::half *)(scales_base + nb * K_SCALE_SIZE); + + 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; + + simd sc = 0; + simd m = 0; + + 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; + + simd qs0 = block_load(qs_base + bi0 * (QK_K / 2)); + simd qs1 = 0; + simd scales0 = block_load(scales_base + bi0 * K_SCALE_SIZE); + simd scales1 = 0; + + const float dall0 = (float)dm_base[bi0 * 2 + 0]; + const float dmin0 = (float)dm_base[bi0 * 2 + 1]; + float dall1 = 0.0f; + float dmin1 = 0.0f; + if (has_row1) { + qs1 = block_load(qs_base + bi1 * (QK_K / 2)); + scales1 = block_load(scales_base + bi1 * K_SCALE_SIZE); + dall1 = (float)dm_base[bi1 * 2 + 0]; + dmin1 = (float)dm_base[bi1 * 2 + 1]; + } + + // unpack Q4_K scale/min codes (get_scale_min_k4 layout), vectorized, per-row + simd scale_f0, min_f0, scale_f1, min_f1; + { + simd scale_lo = scales0.select<4, 1>(0); + simd min_lo = scales0.select<4, 1>(4); + simd hi_bits = scales0.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_f0 = convert(sc) * dall0; + min_f0 = convert(m) * (-dmin0); + } + { + simd scale_lo = scales1.select<4, 1>(0); + simd min_lo = scales1.select<4, 1>(4); + simd hi_bits = scales1.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_f1 = convert(sc) * dall1; + min_f1 = convert(m) * (-dmin1); + } + + simd qs_lo0 = qs0 & simd(0x0f); + simd qs_hi0 = qs0 >> simd(4); + simd qs_lo1 = qs1 & simd(0x0f); + simd qs_hi1 = qs1 >> simd(4); + + // single fused dequant+MAC loop updating both rows each iteration + // the two acc chains are co-scheduled so FMA latency is overlapped + 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 scale0_lo = scale_f0[sb]; + const float scale0_hi = scale_f0[sb + 1]; + const float min0_lo = min_f0[sb]; + const float min0_hi = min_f0[sb + 1]; + const float scale1_lo = scale_f1[sb]; + const float scale1_hi = scale_f1[sb + 1]; + const float min1_lo = min_f1[sb]; + const float min1_hi = min_f1[sb + 1]; + + simd q0_lo = qs_lo0.select<32, 1>(q_offset); + simd q0_hi = qs_hi0.select<32, 1>(q_offset); + simd q1_lo = qs_lo1.select<32, 1>(q_offset); + simd q1_hi = qs_hi1.select<32, 1>(q_offset); + + simd deq0_lo = convert(q0_lo) * scale0_lo + min0_lo; + simd deq0_hi = convert(q0_hi) * scale0_hi + min0_hi; + simd deq1_lo = convert(q1_lo) * scale1_lo + min1_lo; + simd deq1_hi = convert(q1_hi) * scale1_hi + min1_hi; + + acc0 += y_lo * deq0_lo; + acc1 += y_lo * deq1_lo; + acc0 += y_hi * deq0_hi; + acc1 += y_hi * deq1_hi; + } + } + + 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_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) { + // Two partial sums per thread (one per output row of the pair). + 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_q4_k_reorder_esimd( + vx, y, dst, ncols, nrows, lmem, it); + }); + }); +} + +ESIMD_INLINE static void dequantize_mul_mat_vec_q6_k_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; + + const int num_blocks_per_row = ncols / QK_K; + const size_t nb = (size_t)nrows * num_blocks_per_row; + const uint8_t * ql_base = (const uint8_t *)vx; + const uint8_t * qh_base = ql_base + nb * (QK_K / 2); + const int8_t * scales_base = (const int8_t *)(qh_base + nb * (QK_K / 4)); + const sycl::half * d_base = (const sycl::half *)(scales_base + nb * (QK_K / 16)); + + 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; + + simd ql0 = block_load(ql_base + bi0 * (QK_K / 2)); + simd ql1 = 0; + simd qh0 = block_load(qh_base + bi0 * (QK_K / 4)); + simd qh1 = 0; + simd scale0 = block_load(scales_base + bi0 * (QK_K / 16)); + simd scale1 = 0; + if (has_row1) { + ql1 = block_load(ql_base + bi1 * (QK_K / 2)); + qh1 = block_load(qh_base + bi1 * (QK_K / 4)); + scale1 = block_load(scales_base + bi1 * (QK_K / 16)); + } + + simd sc0 = convert(scale0); + simd sc1 = convert(scale1); + const float d0 = (float)d_base[bi0]; + const float d1 = has_row1 ? (float)d_base[bi1] : 0.0f; + + for (int im = 0; im < 2; ++im) { + simd ql_lo0 = ql0.select<32, 1>(64 * im); + simd ql_hi0 = ql0.select<32, 1>(64 * im + 32); + simd qh_bits0 = qh0.select<32, 1>(32 * im); + simd ql_lo1 = ql1.select<32, 1>(64 * im); + simd ql_hi1 = ql1.select<32, 1>(64 * im + 32); + simd qh_bits1 = qh1.select<32, 1>(32 * im); + + // four quant groups, both rows interleaved per group + // reconstruct each 32-wide 6-bit group (matches dequantize_row_q6_K) + for (int g = 0; g < 4; ++g) { + simd y_g = y_vec.select<32, 1>(32 * (4 * im + g)); + + const float scale0_lo = sc0[8 * im + 2 * g + 0] * d0; + const float scale0_hi = sc0[8 * im + 2 * g + 1] * d0; + const float scale1_lo = sc1[8 * im + 2 * g + 0] * d1; + const float scale1_hi = sc1[8 * im + 2 * g + 1] * d1; + + simd scale_vec0; + scale_vec0.select<16, 1>(0) = scale0_lo; + scale_vec0.select<16, 1>(16) = scale0_hi; + simd scale_vec1; + scale_vec1.select<16, 1>(0) = scale1_lo; + scale_vec1.select<16, 1>(16) = scale1_hi; + + simd q0; + simd q1; + switch (g) { + case 0: + q0 = (ql_lo0 & simd(0x0F)) | ((qh_bits0 & simd(0x03)) << simd(4)); + q1 = (ql_lo1 & simd(0x0F)) | ((qh_bits1 & simd(0x03)) << simd(4)); + break; + case 1: + q0 = (ql_hi0 & simd(0x0F)) | ((qh_bits0 & simd(0x0C)) << simd(2)); + q1 = (ql_hi1 & simd(0x0F)) | ((qh_bits1 & simd(0x0C)) << simd(2)); + break; + case 2: + q0 = (ql_lo0 >> simd(4)) | (qh_bits0 & simd(0x30)); + q1 = (ql_lo1 >> simd(4)) | (qh_bits1 & simd(0x30)); + break; + default: + q0 = (ql_hi0 >> simd(4)) | ((qh_bits0 & simd(0xC0)) >> simd(2)); + q1 = (ql_hi1 >> simd(4)) | ((qh_bits1 & simd(0xC0)) >> simd(2)); + break; + } + + simd deq0 = (convert(q0) - 32.0f) * scale_vec0; + simd deq1 = (convert(q1) - 32.0f) * scale_vec1; + + acc0 += y_g * deq0; + acc1 += y_g * deq1; + } + } + } + + 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_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_q6_k_reorder_esimd( + vx, y, dst, ncols, nrows, lmem, it); + }); + }); +} +#endif // GGML_SYCL_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, @@ -1991,7 +2292,11 @@ 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) { +#ifdef GGML_SYCL_ESIMD + dequantize_mul_mat_vec_q4_K_sycl_reorder_esimd(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); +#else dequantize_mul_mat_vec_q4_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); +#endif } else { dequantize_mul_mat_vec_q4_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); } @@ -2007,7 +2312,11 @@ 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) { +#ifdef GGML_SYCL_ESIMD + dequantize_mul_mat_vec_q6_K_sycl_reorder_esimd(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); +#else dequantize_mul_mat_vec_q6_K_sycl_reorder(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); +#endif } else { dequantize_mul_mat_vec_q6_K_sycl(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); } From 3f48e816ac29e61e79080f6e2b39f6ed20e287c0 Mon Sep 17 00:00:00 2001 From: Todd Malsbary Date: Tue, 21 Jul 2026 20:49:27 +0000 Subject: [PATCH 2/8] Refactor ESIMD kernels to share common code Signed-off-by: Todd Malsbary --- ggml/src/ggml-sycl/dmmv.cpp | 247 +++----------------------------- ggml/src/ggml-sycl/esimd.hpp | 269 +++++++++++++++++++++++++++++++++++ 2 files changed, 291 insertions(+), 225 deletions(-) create mode 100644 ggml/src/ggml-sycl/esimd.hpp diff --git a/ggml/src/ggml-sycl/dmmv.cpp b/ggml/src/ggml-sycl/dmmv.cpp index 529912801864..78ffd0188008 100644 --- a/ggml/src/ggml-sycl/dmmv.cpp +++ b/ggml/src/ggml-sycl/dmmv.cpp @@ -15,6 +15,7 @@ #error "GGML_SYCL_ESIMD requires the Intel oneAPI (DPC++) compiler" #endif #include + #include "esimd.hpp" #endif static void convert_f16(const void * vx, const int64_t ib, const int iqs, dfloat2 & v){ @@ -1863,22 +1864,22 @@ static void dequantize_mul_mat_vec_q6_K_sycl(const void *vx, const float *y, } #ifdef GGML_SYCL_ESIMD -constexpr int GGML_SYCL_DMMV_ESIMD_WG_SIZE = 4; - -ESIMD_INLINE static void dequantize_mul_mat_vec_q4_k_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 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 uint8_t * qs_base = (const uint8_t *)vx; - const uint8_t * scales_base = qs_base + nb * (QK_K / 2); - const sycl::half * dm_base = (const sycl::half *)(scales_base + nb * K_SCALE_SIZE); + 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); @@ -1889,96 +1890,13 @@ ESIMD_INLINE static void dequantize_mul_mat_vec_q4_k_reorder_esimd(const void *v simd acc0 = 0.0f; simd acc1 = 0.0f; - simd sc = 0; - simd m = 0; - 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; - - simd qs0 = block_load(qs_base + bi0 * (QK_K / 2)); - simd qs1 = 0; - simd scales0 = block_load(scales_base + bi0 * K_SCALE_SIZE); - simd scales1 = 0; + simd y_vec = block_load(y + (size_t) ib * QK_K); - const float dall0 = (float)dm_base[bi0 * 2 + 0]; - const float dmin0 = (float)dm_base[bi0 * 2 + 1]; - float dall1 = 0.0f; - float dmin1 = 0.0f; - if (has_row1) { - qs1 = block_load(qs_base + bi1 * (QK_K / 2)); - scales1 = block_load(scales_base + bi1 * K_SCALE_SIZE); - dall1 = (float)dm_base[bi1 * 2 + 0]; - dmin1 = (float)dm_base[bi1 * 2 + 1]; - } + 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; - // unpack Q4_K scale/min codes (get_scale_min_k4 layout), vectorized, per-row - simd scale_f0, min_f0, scale_f1, min_f1; - { - simd scale_lo = scales0.select<4, 1>(0); - simd min_lo = scales0.select<4, 1>(4); - simd hi_bits = scales0.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_f0 = convert(sc) * dall0; - min_f0 = convert(m) * (-dmin0); - } - { - simd scale_lo = scales1.select<4, 1>(0); - simd min_lo = scales1.select<4, 1>(4); - simd hi_bits = scales1.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_f1 = convert(sc) * dall1; - min_f1 = convert(m) * (-dmin1); - } - - simd qs_lo0 = qs0 & simd(0x0f); - simd qs_hi0 = qs0 >> simd(4); - simd qs_lo1 = qs1 & simd(0x0f); - simd qs_hi1 = qs1 >> simd(4); - - // single fused dequant+MAC loop updating both rows each iteration - // the two acc chains are co-scheduled so FMA latency is overlapped - 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 scale0_lo = scale_f0[sb]; - const float scale0_hi = scale_f0[sb + 1]; - const float min0_lo = min_f0[sb]; - const float min0_hi = min_f0[sb + 1]; - const float scale1_lo = scale_f1[sb]; - const float scale1_hi = scale_f1[sb + 1]; - const float min1_lo = min_f1[sb]; - const float min1_hi = min_f1[sb + 1]; - - simd q0_lo = qs_lo0.select<32, 1>(q_offset); - simd q0_hi = qs_hi0.select<32, 1>(q_offset); - simd q1_lo = qs_lo1.select<32, 1>(q_offset); - simd q1_hi = qs_hi1.select<32, 1>(q_offset); - - simd deq0_lo = convert(q0_lo) * scale0_lo + min0_lo; - simd deq0_hi = convert(q0_hi) * scale0_hi + min0_hi; - simd deq1_lo = convert(q1_lo) * scale1_lo + min1_lo; - simd deq1_hi = convert(q1_hi) * scale1_hi + min1_hi; - - acc0 += y_lo * deq0_lo; - acc1 += y_lo * deq1_lo; - acc0 += y_hi * deq0_hi; - acc1 += y_hi * deq1_hi; - } + traits::mac_pair(ps, bi0, ps, bi1, has_row1, y_vec, acc0, acc1); } lmem[tid * 2 + 0] = reduce(acc0, std::plus<>{}); @@ -2006,138 +1924,16 @@ static void dequantize_mul_mat_vec_q4_K_sycl_reorder_esimd(const void *vx, const GGML_ASSERT(ncols % QK_K == 0); const int workgroups = (nrows + 1) / 2; stream->submit([&](sycl::handler &h) { - // Two partial sums per thread (one per output row of the pair). 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_q4_k_reorder_esimd( + dequantize_mul_mat_vec_reorder_esimd( vx, y, dst, ncols, nrows, lmem, it); }); }); } -ESIMD_INLINE static void dequantize_mul_mat_vec_q6_k_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; - - const int num_blocks_per_row = ncols / QK_K; - const size_t nb = (size_t)nrows * num_blocks_per_row; - const uint8_t * ql_base = (const uint8_t *)vx; - const uint8_t * qh_base = ql_base + nb * (QK_K / 2); - const int8_t * scales_base = (const int8_t *)(qh_base + nb * (QK_K / 4)); - const sycl::half * d_base = (const sycl::half *)(scales_base + nb * (QK_K / 16)); - - 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; - - simd ql0 = block_load(ql_base + bi0 * (QK_K / 2)); - simd ql1 = 0; - simd qh0 = block_load(qh_base + bi0 * (QK_K / 4)); - simd qh1 = 0; - simd scale0 = block_load(scales_base + bi0 * (QK_K / 16)); - simd scale1 = 0; - if (has_row1) { - ql1 = block_load(ql_base + bi1 * (QK_K / 2)); - qh1 = block_load(qh_base + bi1 * (QK_K / 4)); - scale1 = block_load(scales_base + bi1 * (QK_K / 16)); - } - - simd sc0 = convert(scale0); - simd sc1 = convert(scale1); - const float d0 = (float)d_base[bi0]; - const float d1 = has_row1 ? (float)d_base[bi1] : 0.0f; - - for (int im = 0; im < 2; ++im) { - simd ql_lo0 = ql0.select<32, 1>(64 * im); - simd ql_hi0 = ql0.select<32, 1>(64 * im + 32); - simd qh_bits0 = qh0.select<32, 1>(32 * im); - simd ql_lo1 = ql1.select<32, 1>(64 * im); - simd ql_hi1 = ql1.select<32, 1>(64 * im + 32); - simd qh_bits1 = qh1.select<32, 1>(32 * im); - - // four quant groups, both rows interleaved per group - // reconstruct each 32-wide 6-bit group (matches dequantize_row_q6_K) - for (int g = 0; g < 4; ++g) { - simd y_g = y_vec.select<32, 1>(32 * (4 * im + g)); - - const float scale0_lo = sc0[8 * im + 2 * g + 0] * d0; - const float scale0_hi = sc0[8 * im + 2 * g + 1] * d0; - const float scale1_lo = sc1[8 * im + 2 * g + 0] * d1; - const float scale1_hi = sc1[8 * im + 2 * g + 1] * d1; - - simd scale_vec0; - scale_vec0.select<16, 1>(0) = scale0_lo; - scale_vec0.select<16, 1>(16) = scale0_hi; - simd scale_vec1; - scale_vec1.select<16, 1>(0) = scale1_lo; - scale_vec1.select<16, 1>(16) = scale1_hi; - - simd q0; - simd q1; - switch (g) { - case 0: - q0 = (ql_lo0 & simd(0x0F)) | ((qh_bits0 & simd(0x03)) << simd(4)); - q1 = (ql_lo1 & simd(0x0F)) | ((qh_bits1 & simd(0x03)) << simd(4)); - break; - case 1: - q0 = (ql_hi0 & simd(0x0F)) | ((qh_bits0 & simd(0x0C)) << simd(2)); - q1 = (ql_hi1 & simd(0x0F)) | ((qh_bits1 & simd(0x0C)) << simd(2)); - break; - case 2: - q0 = (ql_lo0 >> simd(4)) | (qh_bits0 & simd(0x30)); - q1 = (ql_lo1 >> simd(4)) | (qh_bits1 & simd(0x30)); - break; - default: - q0 = (ql_hi0 >> simd(4)) | ((qh_bits0 & simd(0xC0)) >> simd(2)); - q1 = (ql_hi1 >> simd(4)) | ((qh_bits1 & simd(0xC0)) >> simd(2)); - break; - } - - simd deq0 = (convert(q0) - 32.0f) * scale_vec0; - simd deq1 = (convert(q1) - 32.0f) * scale_vec1; - - acc0 += y_g * deq0; - acc1 += y_g * deq1; - } - } - } - - 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_q6_K_sycl_reorder_esimd(const void *vx, const float *y, float *dst, const int ncols, const int nrows, @@ -2149,11 +1945,12 @@ static void dequantize_mul_mat_vec_q6_K_sycl_reorder_esimd(const void *vx, const 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_q6_k_reorder_esimd( + dequantize_mul_mat_vec_reorder_esimd( vx, y, dst, ncols, nrows, lmem, it); }); }); } + #endif // GGML_SYCL_ESIMD static void dequantize_mul_mat_vec_q4_K_sycl_reorder(const void *vx, const float *y, diff --git a/ggml/src/ggml-sycl/esimd.hpp b/ggml/src/ggml-sycl/esimd.hpp new file mode 100644 index 000000000000..0bc94ce0bce6 --- /dev/null +++ b/ggml/src/ggml-sycl/esimd.hpp @@ -0,0 +1,269 @@ +// +// 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 + +#if defined(GGML_SYCL_ESIMD) + +#if !defined(__INTEL_LLVM_COMPILER) + #error "GGML_SYCL_ESIMD requires the Intel oneAPI (DPC++) compiler" +#endif + +#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; + +// --------------------------------------------------------------------------- +// 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 }; + } + + // dequantize block bia of pa and block bib of pb, MAC both against y_vec into acc_a / acc_b + // when has_b is false, block b is treated as all-zero and contributes nothing + 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]; + } + + // unpack Q4_K scale/min codes (get_scale_min_k4 layout), vectorized, per block + simd sc = 0; + simd m = 0; + simd scale_f_a, min_f_a, scale_f_b, min_f_b; + { + simd scale_lo = scales_a.select<4, 1>(0); + simd min_lo = scales_a.select<4, 1>(4); + simd hi_bits = scales_a.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_a = convert(sc) * dall_a; + min_f_a = convert(m) * (-dmin_a); + } + { + simd scale_lo = scales_b.select<4, 1>(0); + simd min_lo = scales_b.select<4, 1>(4); + simd hi_bits = scales_b.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_b = convert(sc) * dall_b; + min_f_b = convert(m) * (-dmin_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); + + // single fused dequant+MAC loop updating both accumulators each iteration; + // the two acc chains are co-scheduled so FMA latency is overlapped + 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 scale_a = block_load(pa.scales + bia * (QK_K / 16)); + simd scale_b = 0; + if (has_b) { + ql_b = block_load(pb.ql + bib * (QK_K / 2)); + qh_b = block_load(pb.qh + bib * (QK_K / 4)); + scale_b = block_load(pb.scales + bib * (QK_K / 16)); + } + + simd sc_a = convert(scale_a); + simd sc_b = convert(scale_b); + const float d_a = (float) pa.d[bia]; + const float d_b = has_b ? (float) pb.d[bib] : 0.0f; + + 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); + + // four quant groups, both blocks interleaved per group + // reconstruct each 32-wide 6-bit group (matches dequantize_row_q6_K) + 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; + scale_vec_a.select<16, 1>(0) = scale_a_lo; + scale_vec_a.select<16, 1>(16) = scale_a_hi; + simd scale_vec_b; + scale_vec_b.select<16, 1>(0) = scale_b_lo; + scale_vec_b.select<16, 1>(16) = 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 + +#endif // GGML_SYCL_ESIMD_HPP From 3154b3a63c6d3a3e65718c90a5c812b8de0d40b7 Mon Sep 17 00:00:00 2001 From: Todd Malsbary Date: Tue, 21 Jul 2026 21:51:27 +0000 Subject: [PATCH 3/8] Move control of ESIMD from compile to runtime Signed-off-by: Todd Malsbary --- docs/backend/SYCL.md | 2 +- ggml/CMakeLists.txt | 1 - ggml/src/ggml-sycl/CMakeLists.txt | 4 ---- ggml/src/ggml-sycl/common.hpp | 1 + ggml/src/ggml-sycl/dmmv.cpp | 35 +++++++++++++++++-------------- ggml/src/ggml-sycl/esimd.hpp | 8 ------- ggml/src/ggml-sycl/ggml-sycl.cpp | 8 +++++++ 7 files changed, 29 insertions(+), 30 deletions(-) diff --git a/docs/backend/SYCL.md b/docs/backend/SYCL.md index 934f5bef01db..469a22fd74c3 100644 --- a/docs/backend/SYCL.md +++ b/docs/backend/SYCL.md @@ -776,7 +776,6 @@ use 1 SYCL GPUs: [0] with Max compute units:512 | GGML_SYCL_F16 | OFF *(default)* \|ON *(optional)* | Enable FP16 build with SYCL code path. (1.) | | GGML_SYCL_GRAPH | ON *(default)* \|OFF *(Optional)* | Enable build with [SYCL Graph extension](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc). | | GGML_SYCL_DNN | ON *(default)* \|OFF *(Optional)* | Enable build with oneDNN. | -| GGML_SYCL_ESIMD | OFF *(default)* \|ON *(Optional)* | Enable build with ESIMD kernels. | | GGML_SYCL_HOST_MEM_FALLBACK | ON *(default)* \|OFF *(Optional)* | Allow host memory fallback when device memory is full during quantized weight reorder. Enables inference to continue at reduced speed (reading over PCIe) instead of failing. Requires Linux kernel 6.8+. | | GGML_SYCL_SUPPORT_LEVEL_ZERO_API | ON *(default)* \|OFF *(Optional)* | Support to use Level Zero API for device memory allocation. Requires Level Zero headers/library at build time and Intel GPU driver (Level Zero runtime) at run time. Reduces system RAM usage during multi-GPU inference. SYCL backend always runs on Level Zero running time even if it's set as OFF (The SYCL api will be usage for memory allocation).| | CMAKE_C_COMPILER | `icx` *(Linux)*, `icx/cl` *(Windows)* | Set `icx` compiler for SYCL code path. | @@ -797,6 +796,7 @@ use 1 SYCL GPUs: [0] with Max compute units:512 | GGML_SYCL_ENABLE_DNN | 0 or 1 (default)| Enable running computations through oneDNN and always use oneMKL. | | GGML_SYCL_ENABLE_VMM | 0 or 1 (default) | Enable the virtual-memory device pool. | | 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/CMakeLists.txt b/ggml/CMakeLists.txt index ed059fbaafe0..5381c2136203 100644 --- a/ggml/CMakeLists.txt +++ b/ggml/CMakeLists.txt @@ -251,7 +251,6 @@ option(GGML_SYCL_GRAPH "ggml: enable graphs in the SYCL bac option(GGML_SYCL_HOST_MEM_FALLBACK "ggml: allow host memory fallback in SYCL reorder (requires kernel 6.8+)" ON) option(GGML_SYCL_SUPPORT_LEVEL_ZERO_API "ggml: use Level Zero API in SYCL backend" ON) option(GGML_SYCL_DNN "ggml: enable oneDNN in the SYCL backend" ON) -option(GGML_SYCL_ESIMD "ggml: enable ESIMD kernels in the SYCL backend" OFF) set (GGML_SYCL_TARGET "INTEL" CACHE STRING "ggml: sycl target device") set (GGML_SYCL_DEVICE_ARCH "" CACHE STRING diff --git a/ggml/src/ggml-sycl/CMakeLists.txt b/ggml/src/ggml-sycl/CMakeLists.txt index 332c57c0cf1a..1c17d20df12b 100644 --- a/ggml/src/ggml-sycl/CMakeLists.txt +++ b/ggml/src/ggml-sycl/CMakeLists.txt @@ -158,10 +158,6 @@ else() endif() target_compile_definitions(ggml-sycl PRIVATE GGML_SYCL_DNNL=${GGML_SYCL_DNNL}) -if (GGML_SYCL_ESIMD) - add_compile_definitions(GGML_SYCL_ESIMD) -endif() - if (GGML_SYCL_F16) add_compile_definitions(GGML_SYCL_F16) endif() diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp index e5d9ee89dd86..27d22cc13ae8 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 78ffd0188008..94b8413f86de 100644 --- a/ggml/src/ggml-sycl/dmmv.cpp +++ b/ggml/src/ggml-sycl/dmmv.cpp @@ -8,14 +8,9 @@ #include #define GGML_SYCL_DMMV_HAS_BF16 #endif -#endif - -#if defined(GGML_SYCL_ESIMD) - #if !defined(__INTEL_LLVM_COMPILER) - #error "GGML_SYCL_ESIMD requires the Intel oneAPI (DPC++) compiler" - #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){ @@ -1863,7 +1858,7 @@ static void dequantize_mul_mat_vec_q6_K_sycl(const void *vx, const float *y, }); } -#ifdef GGML_SYCL_ESIMD +#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 @@ -1951,7 +1946,7 @@ static void dequantize_mul_mat_vec_q6_K_sycl_reorder_esimd(const void *vx, const }); } -#endif // GGML_SYCL_ESIMD +#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, @@ -2089,11 +2084,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) { -#ifdef GGML_SYCL_ESIMD - dequantize_mul_mat_vec_q4_K_sycl_reorder_esimd(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); -#else - 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); } @@ -2109,11 +2108,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) { -#ifdef GGML_SYCL_ESIMD - dequantize_mul_mat_vec_q6_K_sycl_reorder_esimd(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); -#else - 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 index 0bc94ce0bce6..838a9a7999e5 100644 --- a/ggml/src/ggml-sycl/esimd.hpp +++ b/ggml/src/ggml-sycl/esimd.hpp @@ -13,12 +13,6 @@ #ifndef GGML_SYCL_ESIMD_HPP #define GGML_SYCL_ESIMD_HPP -#if defined(GGML_SYCL_ESIMD) - -#if !defined(__INTEL_LLVM_COMPILER) - #error "GGML_SYCL_ESIMD requires the Intel oneAPI (DPC++) compiler" -#endif - #include #include "common.hpp" @@ -264,6 +258,4 @@ template <> struct esimd_reorder_q_traits { } // namespace ggml_sycl_esimd -#endif // 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 cb8974eedb75..a0880a3522f5 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -87,6 +87,7 @@ int g_ggml_sycl_enable_dnn = 1; int g_ggml_sycl_fa_onednn = 1; 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; @@ -289,6 +290,7 @@ static void ggml_check_sycl() try { g_ggml_sycl_fa_onednn = ggml_sycl_get_env("GGML_SYCL_FA_ONEDNN", 1); 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); @@ -382,6 +384,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); From d756a6a923bd997335e41c67307c8f9f81d0a038 Mon Sep 17 00:00:00 2001 From: Todd Malsbary Date: Thu, 23 Jul 2026 19:02:12 +0000 Subject: [PATCH 4/8] Use ESIMD by default when available Signed-off-by: Todd Malsbary --- ggml/src/ggml-sycl/ggml-sycl.cpp | 45 +++++++++++++++++++++++--------- 1 file changed, 32 insertions(+), 13 deletions(-) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index a0880a3522f5..f2077871f4ae 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" @@ -3734,6 +3737,21 @@ 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_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: @@ -4437,19 +4455,20 @@ 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 - - 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; - } + // 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 + + if (!g_ggml_sycl_prioritize_dmmv && should_reorder_tensor(ctx, dst)) { + bool use = g_ggml_sycl_enable_esimd && ggml_sycl_supports_reorder_esimd(src0->type); + if (ggml_sycl_supports_reorder_mmvq(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) { From 60d18aaf2aa9febefd7aa1475c017c7080073b9f Mon Sep 17 00:00:00 2001 From: Todd Malsbary Date: Tue, 28 Jul 2026 21:02:03 +0000 Subject: [PATCH 5/8] Fix possible error when using ESIMD by default While not an issue in the current version, this will become an issue when additional QK ESIMD kernels are added (such as Q2_K). Signed-off-by: Todd Malsbary --- ggml/src/ggml-sycl/ggml-sycl.cpp | 24 +++++++++++++----------- 1 file changed, 13 insertions(+), 11 deletions(-) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index f2077871f4ae..ce136fe6ae4d 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -4455,18 +4455,20 @@ 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 - // 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 - - if (!g_ggml_sycl_prioritize_dmmv && should_reorder_tensor(ctx, dst)) { + // 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)))) { bool use = g_ggml_sycl_enable_esimd && ggml_sycl_supports_reorder_esimd(src0->type); - if (ggml_sycl_supports_reorder_mmvq(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; - } + // 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; } From 3cac77ae7a76419779f5ec0681768211579827f8 Mon Sep 17 00:00:00 2001 From: Todd Malsbary Date: Wed, 22 Jul 2026 21:48:04 +0000 Subject: [PATCH 6/8] Add explicit unroll to ESIMD kernels Signed-off-by: Todd Malsbary --- ggml/src/ggml-sycl/esimd.hpp | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-sycl/esimd.hpp b/ggml/src/ggml-sycl/esimd.hpp index 838a9a7999e5..bb0aeb9bfd54 100644 --- a/ggml/src/ggml-sycl/esimd.hpp +++ b/ggml/src/ggml-sycl/esimd.hpp @@ -121,8 +121,9 @@ template <> struct esimd_reorder_q_traits { simd qs_lo_b = qs_b & simd(0x0f); simd qs_hi_b = qs_b >> simd(4); - // single fused dequant+MAC loop updating both accumulators each iteration; + // single fused dequant+MAC loop updating both accumulators each iteration // the two acc chains are co-scheduled so FMA latency is overlapped +#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); @@ -200,6 +201,7 @@ template <> struct esimd_reorder_q_traits { const float d_a = (float) pa.d[bia]; const float d_b = has_b ? (float) pb.d[bib] : 0.0f; +#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); @@ -210,6 +212,7 @@ template <> struct esimd_reorder_q_traits { // four quant groups, both blocks interleaved per group // 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)); From 02d26c7c9c5bd9f5f3d6befa858502f624918d6d Mon Sep 17 00:00:00 2001 From: Todd Malsbary Date: Tue, 28 Jul 2026 21:26:18 +0000 Subject: [PATCH 7/8] Tidy up ESIMD kernels a bit Signed-off-by: Todd Malsbary --- ggml/src/ggml-sycl/esimd.hpp | 117 +++++++++++++++++------------------ 1 file changed, 58 insertions(+), 59 deletions(-) diff --git a/ggml/src/ggml-sycl/esimd.hpp b/ggml/src/ggml-sycl/esimd.hpp index bb0aeb9bfd54..034c3ad137eb 100644 --- a/ggml/src/ggml-sycl/esimd.hpp +++ b/ggml/src/ggml-sycl/esimd.hpp @@ -40,6 +40,39 @@ constexpr int GGML_SYCL_DMMV_ESIMD_WG_SIZE = 4; 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); +} + // --------------------------------------------------------------------------- // 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)] @@ -59,8 +92,6 @@ template <> struct esimd_reorder_q_traits { return { qs, scales, dm }; } - // dequantize block bia of pa and block bib of pb, MAC both against y_vec into acc_a / acc_b - // when has_b is false, block b is treated as all-zero and contributes nothing static ESIMD_INLINE void mac_pair( const ptrs & pa, size_t bia, const ptrs & pb, size_t bib, bool has_b, @@ -85,44 +116,15 @@ template <> struct esimd_reorder_q_traits { dmin_b = (float) pb.dm[bib * 2 + 1]; } - // unpack Q4_K scale/min codes (get_scale_min_k4 layout), vectorized, per block - simd sc = 0; - simd m = 0; simd scale_f_a, min_f_a, scale_f_b, min_f_b; - { - simd scale_lo = scales_a.select<4, 1>(0); - simd min_lo = scales_a.select<4, 1>(4); - simd hi_bits = scales_a.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_a = convert(sc) * dall_a; - min_f_a = convert(m) * (-dmin_a); - } - { - simd scale_lo = scales_b.select<4, 1>(0); - simd min_lo = scales_b.select<4, 1>(4); - simd hi_bits = scales_b.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_b = convert(sc) * dall_b; - min_f_b = convert(m) * (-dmin_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_lo_a = qs_a & simd(0x0F); simd qs_hi_a = qs_a >> simd(4); - simd qs_lo_b = qs_b & simd(0x0f); + simd qs_lo_b = qs_b & simd(0x0F); simd qs_hi_b = qs_b >> simd(4); - // single fused dequant+MAC loop updating both accumulators each iteration - // the two acc chains are co-scheduled so FMA latency is overlapped #pragma unroll for (int sb = 0; sb < 8; sb += 2) { const int q_offset = sb * 16; @@ -169,10 +171,10 @@ template <> struct esimd_reorder_q_traits { }; 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)); + 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 }; } @@ -184,22 +186,24 @@ template <> struct esimd_reorder_q_traits { 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 scale_a = block_load(pa.scales + bia * (QK_K / 16)); - simd scale_b = 0; + 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)); - scale_b = block_load(pb.scales + bib * (QK_K / 16)); + 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(scale_a); - simd sc_b = convert(scale_b); - const float d_a = (float) pa.d[bia]; - const float d_b = has_b ? (float) pb.d[bib] : 0.0f; + simd sc_a = convert(scales_a); + simd sc_b = convert(scales_b); #pragma unroll for (int im = 0; im < 2; ++im) { @@ -210,7 +214,6 @@ template <> struct esimd_reorder_q_traits { simd ql_hi_b = ql_b.select<32, 1>(64 * im + 32); simd qh_bits_b = qh_b.select<32, 1>(32 * im); - // four quant groups, both blocks interleaved per group // reconstruct each 32-wide 6-bit group (matches dequantize_row_q6_K) #pragma unroll for (int g = 0; g < 4; ++g) { @@ -221,12 +224,8 @@ template <> struct esimd_reorder_q_traits { 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; - scale_vec_a.select<16, 1>(0) = scale_a_lo; - scale_vec_a.select<16, 1>(16) = scale_a_hi; - simd scale_vec_b; - scale_vec_b.select<16, 1>(0) = scale_b_lo; - scale_vec_b.select<16, 1>(16) = scale_b_hi; + 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; From 09a9c81394736a43580d241b7aae506bf3223348 Mon Sep 17 00:00:00 2001 From: Todd Malsbary Date: Wed, 12 Aug 2026 17:47:02 +0000 Subject: [PATCH 8/8] Remove redundant copyright notice Signed-off-by: Todd Malsbary --- ggml/src/ggml-sycl/esimd.hpp | 12 ------------ 1 file changed, 12 deletions(-) diff --git a/ggml/src/ggml-sycl/esimd.hpp b/ggml/src/ggml-sycl/esimd.hpp index 034c3ad137eb..cc597e2e3eaf 100644 --- a/ggml/src/ggml-sycl/esimd.hpp +++ b/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