From dcc5f7205351df115720b75c3ec3073198422a3e Mon Sep 17 00:00:00 2001 From: Simon Gardling Date: Fri, 17 Jul 2026 18:20:24 -0400 Subject: [PATCH 1/3] sycl: parallelize the non-contiguous concat kernel Launch geometry only: the non-contiguous concat kernel launched a single-lane work-group (1, 1, 1), now it will launch a (1, 1, SYCL_CONCAT_BLOCK_SIZE) one. SYCL_CONCAT_BLOCK_SIZE is defined in `ggml/src/ggml-sycl/presets.hpp`. llama-bench (Arc Pro B70, Qwen3.6-27B-UD-Q4_K_XL, -fa on, q8_0 KV), on top of upstream master: pp2048 920 -> 1006 t/s (+9.4%) --- ggml/src/ggml-sycl/concat.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-sycl/concat.cpp b/ggml/src/ggml-sycl/concat.cpp index 93e00d65fcd3..0ccd35c24a1f 100644 --- a/ggml/src/ggml-sycl/concat.cpp +++ b/ggml/src/ggml-sycl/concat.cpp @@ -127,7 +127,8 @@ static void concat_T_sycl_non_cont( int64_t ne2, int64_t ne3, uint64_t nb0, uint64_t nb1, uint64_t nb2, uint64_t nb3, int32_t dim) { sycl::range<3> gridDim(ne3, ne2, ne1); - stream->parallel_for(sycl::nd_range<3>(gridDim, sycl::range<3>(1, 1, 1)), [=](sycl::nd_item<3> item_ct1) { + sycl::range<3> blockDim(1, 1, SYCL_CONCAT_BLOCK_SIZE); + stream->parallel_for(sycl::nd_range<3>(gridDim * blockDim, blockDim), [=](sycl::nd_item<3> item_ct1) { int64_t i3 = item_ct1.get_group(0); int64_t i2 = item_ct1.get_group(1); int64_t i1 = item_ct1.get_group(2); From e792674540afb1cdfcb3b246e9cd64d748e7f6a2 Mon Sep 17 00:00:00 2001 From: Simon Gardling Date: Sun, 19 Jul 2026 14:38:02 -0400 Subject: [PATCH 2/3] sycl: cap non-contiguous concat block at ne0 --- ggml/src/ggml-sycl/concat.cpp | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-sycl/concat.cpp b/ggml/src/ggml-sycl/concat.cpp index 0ccd35c24a1f..1ad242fcafbb 100644 --- a/ggml/src/ggml-sycl/concat.cpp +++ b/ggml/src/ggml-sycl/concat.cpp @@ -127,7 +127,14 @@ static void concat_T_sycl_non_cont( int64_t ne2, int64_t ne3, uint64_t nb0, uint64_t nb1, uint64_t nb2, uint64_t nb3, int32_t dim) { sycl::range<3> gridDim(ne3, ne2, ne1); - sycl::range<3> blockDim(1, 1, SYCL_CONCAT_BLOCK_SIZE); + + // Avoid oversubscribing device when there is not enough elements along the innermost dim to + // fill a full SYCL_CONCAT_BLOCK_SIZE. For larger # of elements, the full SYCL_CONCAT_BLOCK_SIZE + // is used. + const int64_t ne0_pad = GGML_PAD(ne0, WARP_SIZE); + const int64_t block_ne0 = ne0_pad < SYCL_CONCAT_BLOCK_SIZE ? ne0_pad : (int64_t) SYCL_CONCAT_BLOCK_SIZE; + sycl::range<3> blockDim(1, 1, block_ne0); + stream->parallel_for(sycl::nd_range<3>(gridDim * blockDim, blockDim), [=](sycl::nd_item<3> item_ct1) { int64_t i3 = item_ct1.get_group(0); int64_t i2 = item_ct1.get_group(1); From 2709909e790d298cc0746950caff16ad8ecf9d6f Mon Sep 17 00:00:00 2001 From: Simon Gardling Date: Mon, 20 Jul 2026 14:46:01 -0400 Subject: [PATCH 3/3] sycl: make non-contiguous concat block width env-tunable (GGML_SYCL_CONCAT_BLOCK_SIZE) --- ggml/src/ggml-sycl/concat.cpp | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-sycl/concat.cpp b/ggml/src/ggml-sycl/concat.cpp index 1ad242fcafbb..5e766a059df8 100644 --- a/ggml/src/ggml-sycl/concat.cpp +++ b/ggml/src/ggml-sycl/concat.cpp @@ -131,8 +131,11 @@ static void concat_T_sycl_non_cont( // Avoid oversubscribing device when there is not enough elements along the innermost dim to // fill a full SYCL_CONCAT_BLOCK_SIZE. For larger # of elements, the full SYCL_CONCAT_BLOCK_SIZE // is used. + // Env-overridable so fewer-core SKUs can try a narrower block without a rebuild. + static const int concat_block_env = ggml_sycl_get_env("GGML_SYCL_CONCAT_BLOCK_SIZE", SYCL_CONCAT_BLOCK_SIZE); + const int64_t concat_block = concat_block_env < WARP_SIZE ? (int64_t) WARP_SIZE : (int64_t) concat_block_env; const int64_t ne0_pad = GGML_PAD(ne0, WARP_SIZE); - const int64_t block_ne0 = ne0_pad < SYCL_CONCAT_BLOCK_SIZE ? ne0_pad : (int64_t) SYCL_CONCAT_BLOCK_SIZE; + const int64_t block_ne0 = ne0_pad < concat_block ? ne0_pad : (int64_t) concat_block; sycl::range<3> blockDim(1, 1, block_ne0); stream->parallel_for(sycl::nd_range<3>(gridDim * blockDim, blockDim), [=](sycl::nd_item<3> item_ct1) {