From cd1bee52038830fb0f943cfd81a3308851a1038f Mon Sep 17 00:00:00 2001 From: cwriter Date: Fri, 2 Oct 2026 10:14:35 +0200 Subject: [PATCH] SYCL: Q8_0 DMMV ESIMD and MMVQ wide load (llama/29186) * Adding wide-load mmvq for Q8_0 and esimd dmmv for q8_0 Assisted-by: Codex * remove guard for q8_0 * remove docs * Simplify by committing to clean code without fallback * Add feature flag as requested Assisted-by: Claude Opus 5 --------- Co-authored-by: cwriter --- ggml/src/ggml-sycl/common.hpp | 1 + ggml/src/ggml-sycl/dmmv.cpp | 149 ++++++++++++++++++++++++++++++- ggml/src/ggml-sycl/ggml-sycl.cpp | 5 +- ggml/src/ggml-sycl/mmvq.cpp | 10 +++ ggml/src/ggml-sycl/vecdotq.hpp | 30 +++++++ 5 files changed, 192 insertions(+), 3 deletions(-) diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp index 661b526a9..11fe86cc0 100644 --- a/ggml/src/ggml-sycl/common.hpp +++ b/ggml/src/ggml-sycl/common.hpp @@ -64,6 +64,7 @@ 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_mmvq_wide; 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 d47d6831a..4f1f97030 100644 --- a/ggml/src/ggml-sycl/dmmv.cpp +++ b/ggml/src/ggml-sycl/dmmv.cpp @@ -2006,6 +2006,135 @@ static void dequantize_mul_mat_vec_q6_K_sycl_reorder_esimd(const void *vx, const }); } +// Q8_0 SOA reorder layout: [qs: nb*QK8_0] [d: nb*sizeof(half)]. +// Process eight blocks per stripe and process remaining blocks one at a time. +template +ESIMD_INLINE void q8_0_mac_stripe( + const int8_t * qs_a, const int8_t * qs_b, + const sycl::half * d_a, const sycl::half * d_b, 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 qa = block_load(qs_a); + simd qb = 0; + // Scale rows can be only 2-byte aligned when nblk_row is odd. + simd da = block_load(d_a, element_aligned_tag{}); + simd db = 0; + if (has_b) { + qb = block_load(qs_b); + db = block_load(d_b, element_aligned_tag{}); + } + + simd da_f = convert(da); + simd db_f = convert(db); + +#pragma unroll + for (int s = 0; s < NBLK; ++s) { + simd y_s = y_vec.template select<32, 1>(s * 32); + simd qa_s = qa.template select<32, 1>(s * 32); + simd qb_s = qb.template select<32, 1>(s * 32); + const float sa = da_f[s]; + const float sb = db_f[s]; + acc_a += y_s * (convert(qa_s) * sa); + acc_b += y_s * (convert(qb_s) * sb); + } +} + +template +ESIMD_INLINE void dequantize_mul_mat_vec_q8_0_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; + + constexpr int STRIPE = 8; + + const int nblk_row = ncols / QK8_0; + const size_t nb = (size_t) nrows * nblk_row; + const int8_t * qs = (const int8_t *) vx; + const sycl::half * d = (const sycl::half *) (qs + nb * QK8_0); + + const int tid = it.get_local_id(0); + const int row_pair = it.get_group(0); + const int row0 = row_pair * 2; + const bool has_row1 = row0 + 1 < nrows; + + const size_t base0 = (size_t) row0 * nblk_row; + const size_t base1 = has_row1 ? (size_t) (row0 + 1) * nblk_row : base0; + + simd acc0 = 0.0f; + simd acc1 = 0.0f; + + // Each thread processes one contiguous stripe. + int ib = 0; + for (; ib + WG * STRIPE <= nblk_row; ib += WG * STRIPE) { + const int b = ib + tid * STRIPE; + simd y_vec = block_load(y + (size_t) b * QK8_0); + q8_0_mac_stripe(qs + (base0 + b) * QK8_0, qs + (base1 + b) * QK8_0, + d + base0 + b, d + base1 + b, has_row1, y_vec, acc0, acc1); + } + + // Distribute remaining blocks across the work-group. + for (int b = ib + tid; b < nblk_row; b += WG) { + simd y_vec = block_load(y + (size_t) b * QK8_0); + q8_0_mac_stripe<1>(qs + (base0 + b) * QK8_0, qs + (base1 + b) * QK8_0, + d + base0 + b, d + base1 + b, 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 < WG; ++p) { + sum0 += lmem[p * 2 + 0]; + sum1 += lmem[p * 2 + 1]; + } + dst[row0 + 0] = sum0; + if (has_row1) { + dst[row0 + 1] = sum1; + } + } +} + +template +static void q8_0_esimd_launch(const void * vx, const float * y, float * dst, const int ncols, + const int nrows, dpct::queue_ptr stream) { + const int workgroups = (nrows + 1) / 2; + stream->submit([&](sycl::handler & h) { + sycl::local_accessor lmem(sycl::range<1>(WG * 2), h); + h.parallel_for( + sycl::nd_range<1>(sycl::range<1>((size_t) workgroups * WG), sycl::range<1>(WG)), + [=](sycl::nd_item<1> it) [[intel::sycl_explicit_simd]] { + dequantize_mul_mat_vec_q8_0_reorder_esimd(vx, y, dst, ncols, nrows, lmem, it); + }); + }); +} + +static void dequantize_mul_mat_vec_q8_0_sycl_reorder_esimd(const void *vx, const float *y, + float *dst, const int ncols, + const int nrows, + dpct::queue_ptr stream) { + GGML_ASSERT(ncols % QK8_0 == 0); + + // Scale the work-group with the number of blocks per row. + const int nblk_row = ncols / QK8_0; + if (nblk_row >= 64) { + q8_0_esimd_launch<8>(vx, y, dst, ncols, nrows, stream); + } else if (nblk_row >= 32) { + q8_0_esimd_launch<4>(vx, y, dst, ncols, nrows, stream); + } else if (nblk_row >= 16) { + q8_0_esimd_launch<2>(vx, y, dst, ncols, nrows, stream); + } else { + q8_0_esimd_launch<1>(vx, y, dst, ncols, nrows, stream); + } +} + #endif // GGML_SYCL_DMMV_HAS_ESIMD static void dequantize_mul_mat_vec_q4_K_sycl_reorder(const void *vx, const float *y, @@ -2072,11 +2201,20 @@ void ggml_sycl_op_dequantize_mul_mat_vec( ggml_sycl_pool_alloc src1_dfloat_a(ctx.pool()); sycl::half *src1_dfloat = nullptr; // dfloat == half +#ifdef GGML_SYCL_DMMV_HAS_ESIMD + // The ESIMD Q8_0 kernel reads F32 activations. + const bool q8_0_esimd = src0->type == GGML_TYPE_Q8_0 && g_ggml_sycl_enable_esimd && + ((ggml_tensor_extra_gpu *) dst->src[0]->extra) && + ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder; +#else + const bool q8_0_esimd = false; +#endif + bool src1_convert_f16 = src0->type == GGML_TYPE_Q1_0 || src0->type == GGML_TYPE_Q4_0 || src0->type == GGML_TYPE_Q4_1 || src0->type == GGML_TYPE_Q5_0 || src0->type == GGML_TYPE_Q5_1 || - src0->type == GGML_TYPE_Q8_0 || src0->type == GGML_TYPE_F16 || + (src0->type == GGML_TYPE_Q8_0 && !q8_0_esimd) || src0->type == GGML_TYPE_F16 || src0->type == GGML_TYPE_BF16; if (src1_convert_f16) { @@ -2120,7 +2258,14 @@ void ggml_sycl_op_dequantize_mul_mat_vec( case GGML_TYPE_Q8_0: if ((ggml_tensor_extra_gpu *) dst->src[0]->extra && ((ggml_tensor_extra_gpu *) dst->src[0]->extra)->optimized_feature.reorder) { - dequantize_mul_mat_vec_q8_0_sycl_reorder(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream); +#ifdef GGML_SYCL_DMMV_HAS_ESIMD + if (g_ggml_sycl_enable_esimd) { + dequantize_mul_mat_vec_q8_0_sycl_reorder_esimd(src0_dd_i, src1_ddf_i, dst_dd_i, ne00, row_diff, stream); + } else +#endif + { + dequantize_mul_mat_vec_q8_0_sycl_reorder(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream); + } } else { dequantize_mul_mat_vec_q8_0_sycl(src0_dd_i, src1_dfloat, dst_dd_i, ne00, row_diff, stream); } diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index e4e61e306..49148d7c5 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -103,6 +103,7 @@ int g_ggml_sycl_memtrace_step = 64; 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_mmvq_wide = 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; @@ -398,6 +399,7 @@ static void ggml_check_sycl() try { 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_mmvq_wide = ggml_sycl_get_env("GGML_SYCL_MMVQ_WIDE", 1); g_ggml_sycl_prioritize_dmmv = ggml_sycl_get_env("GGML_SYCL_PRIORITIZE_DMMV", 0); #ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API @@ -521,7 +523,7 @@ static void ggml_check_sycl() try { #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_MMVQ_WIDE: %d\n", g_ggml_sycl_mmvq_wide); 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); @@ -4142,6 +4144,7 @@ static bool ggml_sycl_supports_reorder_esimd(enum ggml_type type) { case GGML_TYPE_Q4_K: case GGML_TYPE_Q5_K: case GGML_TYPE_Q6_K: + case GGML_TYPE_Q8_0: return true; default: return false; diff --git a/ggml/src/ggml-sycl/mmvq.cpp b/ggml/src/ggml-sycl/mmvq.cpp index 7e4f22dd1..d0f090703 100644 --- a/ggml/src/ggml-sycl/mmvq.cpp +++ b/ggml/src/ggml-sycl/mmvq.cpp @@ -1192,6 +1192,16 @@ static void reorder_mul_mat_vec_q8_0_q8_1_sycl(const void * vx, const void * vy, const sycl::range<3> block_nums(1, 1, block_num_y); const sycl::range<3> block_dims(1, GGML_SYCL_MMV_Y, num_subgroups * WARP_SIZE); + if (g_ggml_sycl_mmvq_wide) { + stream->submit([&](sycl::handler & cgh) { + cgh.parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), + [=](sycl::nd_item<3> nd_item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] { + mul_mat_vec_q_reorder(vx, vy, dst, ncols, nrows, nd_item); + }); + }); + return; + } + stream->submit([&](sycl::handler & cgh) { cgh.parallel_for(sycl::nd_range<3>(block_nums * block_dims, block_dims), [=](sycl::nd_item<3> nd_item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] { diff --git a/ggml/src/ggml-sycl/vecdotq.hpp b/ggml/src/ggml-sycl/vecdotq.hpp index 909f7a789..cc6ae6a8d 100644 --- a/ggml/src/ggml-sycl/vecdotq.hpp +++ b/ggml/src/ggml-sycl/vecdotq.hpp @@ -415,6 +415,36 @@ template <> struct reorder_vec_dot_q_sycl { }; }; +// Load four contiguous dwords per operand instead of loading each value separately. +struct reorder_vec_dot_q8_0_wide { + static constexpr ggml_type gtype = GGML_TYPE_Q8_0; + + using q8_0_block = ggml_sycl_reordered::block_q_t; + using q8_0_traits = typename q8_0_block::traits; + + __dpct_inline__ float operator()(const void * __restrict__ vbq, const std::pair ibx_offset, + const std::pair d_offset, const int8_t * q8_1_quant_ptr, + const sycl::half2 * q8_1_ds, const int & iqs) { + static_assert(q8_0_traits::vdr_mmvq == 4, "the wide load moves exactly four dwords"); + + const uint8_t * base = static_cast(vbq); + const int8_t * qs = reinterpret_cast(base + ibx_offset.first); + const ggml_half d = *reinterpret_cast(base + d_offset.first); + + const sycl::int4 v = *reinterpret_cast(qs + sizeof(int) * iqs); + const sycl::int4 u = *reinterpret_cast(q8_1_quant_ptr + sizeof(int) * iqs); + + int sumi = 0; +#pragma unroll + for (int i = 0; i < 4; ++i) { + sumi = dpct::dp4a(v[i], u[i], sumi); + } + + const sycl::half2 ds_values = *q8_1_ds; + return static_cast(d) * static_cast(ds_values[0]) * sumi; + } +}; + template <> struct reorder_vec_dot_q_sycl { static constexpr ggml_type gtype = GGML_TYPE_Q8_0;