Skip to content

sycl: parallelize the non-contiguous concat kernel #25852

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft
Titaniumtown wants to merge 3 commits into
base: master
Choose a base branch
from
+12 −1
Draft

sycl: parallelize the non-contiguous concat kernel #25852

Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
dcc5f72
sycl: parallelize the non-contiguous concat kernel
Titaniumtown Jul 17, 2026
e792674
sycl: cap non-contiguous concat block at ne0
Titaniumtown Jul 19, 2026
2709909
sycl: make non-contiguous concat block width env-tunable (GGML_SYCL_C…
Titaniumtown Jul 20, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
13 changes: 12 additions & 1 deletion ggml/src/ggml-sycl/concat.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -127,7 +127,18 @@ 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) {

// 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 < 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) {
int64_t i3 = item_ct1.get_group(0);
int64_t i2 = item_ct1.get_group(1);
int64_t i1 = item_ct1.get_group(2);
Expand Down