From 99b36ec76585d090759839aeb9b132c64ae504f4 Mon Sep 17 00:00:00 2001 From: Neo Zhang Jianyu Date: Fri, 18 Jul 2025 10:23:14 +0800 Subject: [PATCH 1/2] use max work group size for device to replace the magic number (#14732) --- ggml/src/ggml-sycl/ggml-sycl.cpp | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 128b6d56040..8058bbb7cac 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -3010,8 +3010,11 @@ static void ggml_sycl_mul_mat_id(ggml_backend_sycl_context & ctx, SYCL_CHECK(CHECK_TRY_ERROR( stream->memset(dev_cur_src1_row.get(), 0, sizeof(int)))); + const unsigned int max_work_group_size = ggml_sycl_info().max_work_group_sizes[ctx.device]; + assert(work_group_size % (WARP_SIZE * WARP_SIZE) == 0); + { - sycl::range<3> block_dims(1, 1, std::min((unsigned int)ne10, 768u)); + sycl::range<3> block_dims(1, 1, std::min((unsigned int)ne10, max_work_group_size)); sycl::range<3> grid_dims(1, n_ids, ids->ne[1]); stream->submit([&](sycl::handler &cgh) { sycl::local_accessor src1_row_acc(cgh); @@ -3056,7 +3059,7 @@ static void ggml_sycl_mul_mat_id(ggml_backend_sycl_context & ctx, ggml_sycl_mul_mat(ctx, &src0_row, &src1_row, &dst_row); { - sycl::range<3> block_dims(1, 1, std::min((unsigned int)ne0, 768u)); + sycl::range<3> block_dims(1, 1, std::min((unsigned int)ne0, max_work_group_size)); sycl::range<3> grid_dims(1, 1, num_src1_rows); stream->submit([&](sycl::handler &cgh) { const char *__restrict dst_contiguous_get = From e37b33d4a4de44fb02406ea198345a9e85885dba Mon Sep 17 00:00:00 2001 From: arthw <14088817+arthw@users.noreply.github.com> Date: Mon, 21 Jul 2025 17:51:23 +0800 Subject: [PATCH 2/2] update for local api --- ggml/src/ggml-sycl/ggml-sycl.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 8058bbb7cac..24da7c3c9a4 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -3010,7 +3010,7 @@ static void ggml_sycl_mul_mat_id(ggml_backend_sycl_context & ctx, SYCL_CHECK(CHECK_TRY_ERROR( stream->memset(dev_cur_src1_row.get(), 0, sizeof(int)))); - const unsigned int max_work_group_size = ggml_sycl_info().max_work_group_sizes[ctx.device]; + const unsigned int max_work_group_size = ggml_sycl_info().work_group_size(ctx.device); assert(work_group_size % (WARP_SIZE * WARP_SIZE) == 0); {