From d209086157f007ffa672991518634b8f283522b2 Mon Sep 17 00:00:00 2001 From: Todd Malsbary Date: Tue, 7 Jul 2026 00:43:41 -0700 Subject: [PATCH] sycl : set K_QUANTS_PER_ITERATION to 1 on DMMV path (#25063) * sycl: add supported types to ggml_sycl_supports_reorder_dmmv The reordered feature is implemented in ggml_sycl_op_dequantize_mul_mat_vec, but gated by ggml_sycl_supports_reorder_dmmv. This commit fixes the gate. Signed-off-by: Todd Malsbary * sycl: set K_QUANTS_PER_ITERATION=1 to improve utilization When combined with opening the reorder gate, this improves GPU utilization on B70, giving a significant boost to tg t/s. Signed-off-by: Todd Malsbary * sycl: replace QK_WARP_SIZE with WARP_SIZE for QK_5 Signed-off-by: Todd Malsbary * sycl: add missing types to ggml_backend_sycl_buffer_init_tensor Without this, the extra field is not allocated and the reorder path will not take effect. Signed-off-by: Todd Malsbary --------- Signed-off-by: Todd Malsbary --- ggml/src/ggml-sycl/dmmv.cpp | 27 +++++++++++++++------------ ggml/src/ggml-sycl/ggml-sycl.cpp | 6 ++++++ ggml/src/ggml-sycl/presets.hpp | 2 +- 3 files changed, 22 insertions(+), 13 deletions(-) diff --git a/ggml/src/ggml-sycl/dmmv.cpp b/ggml/src/ggml-sycl/dmmv.cpp index fb8a1757f1..5c6835e1d2 100644 --- a/ggml/src/ggml-sycl/dmmv.cpp +++ b/ggml/src/ggml-sycl/dmmv.cpp @@ -680,14 +680,14 @@ static void dequantize_mul_mat_vec_q4_k(const void *__restrict__ vx, q16[2] = q2[0] & 0x0f0f; q16[3] = q2[0] & 0xf0f0; - float4 s = {0.f, 0.f, 0.f, 0.f}; + sycl::float4 s = {0.f, 0.f, 0.f, 0.f}; float smin = 0; for (int l = 0; l < 2; ++l) { - s.x += y1[l] * q4[l+0]; s.y += y1[l+32] * q4[l+2]; - s.z += y2[l] * q4[l+4]; s.w += y2[l+32] * q4[l+6]; + s.x() += y1[l] * q4[l+0]; s.y() += y1[l+32] * q4[l+2]; + s.z() += y2[l] * q4[l+4]; s.w() += y2[l+32] * q4[l+6]; smin += y1[l] * sc[2] + y1[l+32] * sc[3] + y2[l] * sc[6] + y2[l+32] * sc[7]; } - tmp += dall * (s.x * sc[0] + s.y * sc[1] * 1.f/16.f + s.z * sc[4] + s.w * sc[5] * 1.f/16.f) - dmin * smin; + tmp += dall * (s.x() * sc[0] + s.y() * sc[1] * 1.f/16.f + s.z() * sc[4] + s.w() * sc[5] * 1.f/16.f) - dmin * smin; #endif } @@ -835,14 +835,14 @@ static void dequantize_mul_mat_vec_q4_k_reorder(const void *__restrict__ vx, q16[2] = q2[0] & 0x0f0f; q16[3] = q2[0] & 0xf0f0; - float4 s = {0.f, 0.f, 0.f, 0.f}; + sycl::float4 s = {0.f, 0.f, 0.f, 0.f}; float smin = 0; for (int l = 0; l < 2; ++l) { - s.x += y1[l] * q4[l+0]; s.y += y1[l+32] * q4[l+2]; - s.z += y2[l] * q4[l+4]; s.w += y2[l+32] * q4[l+6]; + s.x() += y1[l] * q4[l+0]; s.y() += y1[l+32] * q4[l+2]; + s.z() += y2[l] * q4[l+4]; s.w() += y2[l+32] * q4[l+6]; smin += y1[l] * sc[2] + y1[l+32] * sc[3] + y2[l] * sc[6] + y2[l+32] * sc[7]; } - tmp += dall * (s.x * sc[0] + s.y * sc[1] * 1.f/16.f + s.z * sc[4] + s.w * sc[5] * 1.f/16.f) - dmin * smin; + tmp += dall * (s.x() * sc[0] + s.y() * sc[1] * 1.f/16.f + s.z() * sc[4] + s.w() * sc[5] * 1.f/16.f) - dmin * smin; #endif } @@ -1126,7 +1126,7 @@ static void dequantize_mul_mat_vec_q5_k_reorder(const void *__restrict__ vx, // sum up partial sums and write back result #pragma unroll - for (int mask = QK_WARP_SIZE / 2; mask > 0; mask >>= 1) { + for (int mask = WARP_SIZE / 2; mask > 0; mask >>= 1) { tmp += dpct::permute_sub_group_by_xor(item_ct1.get_sub_group(), tmp, mask); } @@ -1762,10 +1762,13 @@ static void dequantize_mul_mat_vec_q5_K_sycl_reorder(const void *vx, const float const int nrows, dpct::queue_ptr stream) { GGML_ASSERT(ncols % QK_K == 0); - const sycl::range<3> block_dims(1, 1, QK_WARP_SIZE); + const int ny = 2 / K_QUANTS_PER_ITERATION; + const int block_num_y = (nrows + ny - 1) / ny; + const sycl::range<3> block_nums(1, 1, block_num_y); + const sycl::range<3> block_dims(1, ny, WARP_SIZE); stream->parallel_for( - sycl::nd_range<3>(sycl::range<3>(1, 1, nrows) * block_dims, block_dims), - [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(QK_WARP_SIZE)]] { + sycl::nd_range<3>(block_nums * block_dims, block_dims), + [=](sycl::nd_item<3> item_ct1) [[sycl::reqd_sub_group_size(WARP_SIZE)]] { dequantize_mul_mat_vec_q5_k_reorder(vx, y, dst, ncols, nrows, item_ct1); }); } diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 0b7576ca3f..48256e2b53 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -546,7 +546,9 @@ ggml_backend_sycl_buffer_init_tensor(ggml_backend_buffer_t buffer, switch (tensor->type) { case GGML_TYPE_Q4_0: case GGML_TYPE_Q8_0: + case GGML_TYPE_Q3_K: case GGML_TYPE_Q4_K: + case GGML_TYPE_Q5_K: case GGML_TYPE_Q6_K:{ ggml_tensor_extra_gpu * extra = new ggml_tensor_extra_gpu{}; tensor->extra = extra; @@ -3687,6 +3689,10 @@ inline bool ggml_sycl_supports_reorder_dmmv(enum ggml_type type) { case GGML_TYPE_Q1_0: case GGML_TYPE_Q4_0: case GGML_TYPE_Q8_0: + case GGML_TYPE_Q3_K: + case GGML_TYPE_Q4_K: + case GGML_TYPE_Q5_K: + case GGML_TYPE_Q6_K: return true; default: return false; diff --git a/ggml/src/ggml-sycl/presets.hpp b/ggml/src/ggml-sycl/presets.hpp index 54566316ad..3bd1c86eae 100644 --- a/ggml/src/ggml-sycl/presets.hpp +++ b/ggml/src/ggml-sycl/presets.hpp @@ -62,7 +62,7 @@ #endif #ifndef K_QUANTS_PER_ITERATION -#define K_QUANTS_PER_ITERATION 2 +#define K_QUANTS_PER_ITERATION 1 #else static_assert(K_QUANTS_PER_ITERATION == 1 || K_QUANTS_PER_ITERATION == 2, "K_QUANTS_PER_ITERATION must be 1 or 2"); #endif