upstream-archive/llama.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/llama.cpp.git synced 2026-08-06 10:00:48 -05:00
Code Issues Packages Projects Releases Wiki Activity

sycl: parallelize the non-contiguous concat kernel (#25852)

Browse Source 
* 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%)

* sycl: cap non-contiguous concat block at ne0

* sycl: make non-contiguous concat block width env-tunable (GGML_SYCL_CONCAT_BLOCK_SIZE)

* Revert "sycl: make non-contiguous concat block width env-tunable (GGML_SYCL_CONCAT_BLOCK_SIZE)"

This reverts commit 2709909e79.
This commit is contained in:
Titaniumtown
2026-08-03 19:08:05 -07:00
committed by GitHub
parent 66fa168a56
commit 6c8dcaa7ae
1 changed files with 9 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

10
ggml/src/ggml-sycl/concat.cpp
View File

@@ -127,7 +127,15 @@ 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.
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);
int64_t i1 = item_ct1.get_group(2);