SYCL: Add non contiguous support in RMS_NORM and NORM kernels #13611

qnixsynapse · 2025-05-18T06:39:49Z

Added non contiguous support in RMS_NORM and NORM kernels. test-backend-ops seems to pass with this change.

Edit: restored logic for handling multi subgroup correctly which was not tested by test-backend-ops

qnixsynapse · 2025-05-19T06:16:09Z

It now seems to pass with ne[0] = 1920

 NORM(type=f32,ne=[64,5,4,3],v=0,eps=0.000000): OK
  NORM(type=f32,ne=[1920,5,4,3],v=0,eps=0.000000): OK
  NORM(type=f32,ne=[64,5,4,3],v=1,eps=0.000000): OK
  NORM(type=f32,ne=[1920,5,4,3],v=1,eps=0.000000): OK
  NORM(type=f32,ne=[64,5,4,3],v=0,eps=0.000001): OK
  NORM(type=f32,ne=[1920,5,4,3],v=0,eps=0.000001): OK
  NORM(type=f32,ne=[64,5,4,3],v=1,eps=0.000001): OK
  NORM(type=f32,ne=[1920,5,4,3],v=1,eps=0.000001): OK
  NORM(type=f32,ne=[64,5,4,3],v=0,eps=0.000100): OK
  NORM(type=f32,ne=[1920,5,4,3],v=0,eps=0.000100): OK
  NORM(type=f32,ne=[64,5,4,3],v=1,eps=0.000100): OK
  NORM(type=f32,ne=[1920,5,4,3],v=1,eps=0.000100): OK
  NORM(type=f32,ne=[64,5,4,3],v=0,eps=0.100000): OK
  NORM(type=f32,ne=[1920,5,4,3],v=0,eps=0.100000): OK
  NORM(type=f32,ne=[64,5,4,3],v=1,eps=0.100000): OK
  NORM(type=f32,ne=[1920,5,4,3],v=1,eps=0.100000): OK
  RMS_NORM(type=f32,ne=[64,5,4,3],v=0,eps=0.000000): OK
  RMS_NORM(type=f32,ne=[1920,5,4,3],v=0,eps=0.000000): OK
  RMS_NORM(type=f32,ne=[64,5,4,3],v=1,eps=0.000000): OK
  RMS_NORM(type=f32,ne=[1920,5,4,3],v=1,eps=0.000000): OK
  RMS_NORM(type=f32,ne=[64,5,4,3],v=0,eps=0.000001): OK
  RMS_NORM(type=f32,ne=[1920,5,4,3],v=0,eps=0.000001): OK
  RMS_NORM(type=f32,ne=[64,5,4,3],v=1,eps=0.000001): OK
  RMS_NORM(type=f32,ne=[1920,5,4,3],v=1,eps=0.000001): OK
  RMS_NORM(type=f32,ne=[64,5,4,3],v=0,eps=0.000100): OK
  RMS_NORM(type=f32,ne=[1920,5,4,3],v=0,eps=0.000100): OK
  RMS_NORM(type=f32,ne=[64,5,4,3],v=1,eps=0.000100): OK
  RMS_NORM(type=f32,ne=[1920,5,4,3],v=1,eps=0.000100): OK
  RMS_NORM(type=f32,ne=[64,5,4,3],v=0,eps=0.100000): OK
  RMS_NORM(type=f32,ne=[1920,5,4,3],v=0,eps=0.100000): OK
  RMS_NORM(type=f32,ne=[64,5,4,3],v=1,eps=0.100000): OK
  RMS_NORM(type=f32,ne=[1920,5,4,3],v=1,eps=0.100000): OK
  5535/5535 tests passed
  Backend SYCL0: OK

Rbiessy

Were you able to measure the impact on performance for this change? If it has one we may want to introduce different paths for contiguous and non-contiguous cases.

I'll try to check on my side for some relevant sizes at some point.

qnixsynapse · 2025-05-19T11:50:14Z

Were you able to measure the impact on performance for this change? If it has one we may want to introduce different paths for contiguous and non-contiguous cases.

I'll try to check on my side for some relevant sizes at some point.

Doesn't seem much different from master:

[llama.cpp][master]$ build/bin/llama-bench -ngl 99 -m ~/Downloads/Weights/pythia-1.4b-q4_0.gguf

model	size	params	backend	ngl	test	t/s
gptneox 1.4B Q4_0	786.31 MiB	1.41 B	SYCL	99	pp512	4202.02 ± 7.96
gptneox 1.4B Q4_0	786.31 MiB	1.41 B	SYCL	99	tg128	46.34 ± 0.25

build: 92ecdcc (5423)

[llama.cpp][sycl/non_cont_norms]$ build/bin/llama-bench -ngl 99 -m ~/Downloads/Weights/pythia-1.4b-q4_0.gguf

model	size	params	backend	ngl	test	t/s
gptneox 1.4B Q4_0	786.31 MiB	1.41 B	SYCL	99	pp512	4202.28 ± 8.47
gptneox 1.4B Q4_0	786.31 MiB	1.41 B	SYCL	99	tg128	46.25 ± 0.26

build: This PR

Alcpz · 2025-05-20T13:34:18Z

I ran a couple of models and don't see any regressions there. It seems the changes don't affect the existing codepath (e2e):

model	device	test	avg_ts (b5424)	avg_ts (b5429)	speedup
qwen2 1.5B Q4_0	Arc A770	pp512	4042.61	4020.68	0.99
qwen2 1.5B Q4_0	Arc A770	tg128	46.9076	46.919	1
qwen2 1.5B Q4_K - Medium	Arc A770	pp512	4418.48	4408.21	1
qwen2 1.5B Q4_K - Medium	Arc A770	tg128	46.3589	45.9716	0.99
llama 7B Q4_0	Arc A770	pp512	1460.58	1460.19	1
llama 7B Q4_0	Arc A770	tg128	34.986	34.7013	0.99
llama 7B Q4_K - Medium	Arc A770	pp512	1718.93	1717.47	1
llama 7B Q4_K - Medium	Arc A770	tg128	32.7239	32.8426	1
gemma2 2B Q4_K - Medium	Arc A770	pp512	3584.07	3581.23	1
gemma2 2B Q4_K - Medium	Arc A770	tg128	39.4594	39.3022	1
phi3 3B Q4_0	Arc A770	pp512	2153.61	2156.82	1
phi3 3B Q4_0	Arc A770	tg128	40.3495	40.2226	1
phi3 3B Q4_K - Medium	Arc A770	pp512	2492.66	2491.55	1
phi3 3B Q4_K - Medium	Arc A770	tg128	34.7158	34.788	1
llama 34B Q6_K	Arc A770	pp512	1023.89	1021.92	1
llama 34B Q6_K	Arc A770	tg128	15.2625	15.2357	1
llama 8B Q4_K - Medium	Arc A770	pp512	1615.01	1616.76	1
llama 8B Q4_K - Medium	Arc A770	tg128	28.8256	28.7515	1

AD2605 · 2025-05-20T14:06:05Z

ggml/src/ggml-sycl/norm.cpp

+    if  (block_size > WARP_SIZE) {
+        int warp_id = tid / WARP_SIZE;
+        int lane_id = tid % WARP_SIZE;


I know this was already in the existing code, but in case you'd like to change it to a more sycl like fashion, you can remove the if (block_size > WARP_SIZE) and do

auto sub_group = item_ct1.get_sub_group(); auto sg_id = sub_group.get_group_linear_id(); auto wi_in_sg = sub_group.get_local_linear_id(); // And this is same as lane_id

You can refer https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sub-group-class

I don't think this will work with every ne[0] sizes:

NORM(type=f32,ne=[64,5,4,3],v=0,eps=0.000000): [NORM] inf mismatch: SYCL0=inf CPU=0.103443 FAIL NORM(type=f32,ne=[1920,5,4,3],v=0,eps=0.000000): OK NORM(type=f32,ne=[64,5,4,3],v=1,eps=0.000000): [NORM] inf mismatch: SYCL0=inf CPU=1.217545 FAIL NORM(type=f32,ne=[1920,5,4,3],v=1,eps=0.000000): [NORM] inf mismatch: SYCL0=-inf CPU=-0.915058 FAIL NORM(type=f32,ne=[64,5,4,3],v=0,eps=0.000001): [NORM] NMSE = 0.996613428 > 0.000000100 FAIL NORM(type=f32,ne=[1920,5,4,3],v=0,eps=0.000001): OK NORM(type=f32,ne=[64,5,4,3],v=1,eps=0.000001): [NORM] NMSE = 0.996409533 > 0.000000100 FAIL NORM(type=f32,ne=[1920,5,4,3],v=1,eps=0.000001): [NORM] NMSE = 0.996592679 > 0.000000100 FAIL NORM(type=f32,ne=[64,5,4,3],v=0,eps=0.000100): [NORM] NMSE = 0.966021524 > 0.000000100 FAIL NORM(type=f32,ne=[1920,5,4,3],v=0,eps=0.000100): OK NORM(type=f32,ne=[64,5,4,3],v=1,eps=0.000100): [NORM] NMSE = 0.965700277 > 0.000000100 FAIL NORM(type=f32,ne=[1920,5,4,3],v=1,eps=0.000100): [NORM] NMSE = 0.965806282 > 0.000000100 FAIL NORM(type=f32,ne=[64,5,4,3],v=0,eps=0.100000): [NORM] NMSE = 0.282048698 > 0.000000100 FAIL NORM(type=f32,ne=[1920,5,4,3],v=0,eps=0.100000): OK NORM(type=f32,ne=[64,5,4,3],v=1,eps=0.100000): [NORM] NMSE = 0.285614294 > 0.000000100 FAIL NORM(type=f32,ne=[1920,5,4,3],v=1,eps=0.100000): [NORM] NMSE = 0.270144864 > 0.000000100 FAIL RMS_NORM(type=f32,ne=[64,5,4,3],v=0,eps=0.000000): [RMS_NORM] inf mismatch: SYCL0=-inf CPU=-0.151891 FAIL RMS_NORM(type=f32,ne=[1920,5,4,3],v=0,eps=0.000000): OK RMS_NORM(type=f32,ne=[64,5,4,3],v=1,eps=0.000000): [RMS_NORM] inf mismatch: SYCL0=inf CPU=1.350047 FAIL RMS_NORM(type=f32,ne=[1920,5,4,3],v=1,eps=0.000000): [RMS_NORM] inf mismatch: SYCL0=inf CPU=1.245914 FAIL RMS_NORM(type=f32,ne=[64,5,4,3],v=0,eps=0.000001): [RMS_NORM] NMSE = 0.999657826 > 0.000000100 FAIL RMS_NORM(type=f32,ne=[1920,5,4,3],v=0,eps=0.000001): OK RMS_NORM(type=f32,ne=[64,5,4,3],v=1,eps=0.000001): [RMS_NORM] NMSE = 0.999642137 > 0.000000100 FAIL RMS_NORM(type=f32,ne=[1920,5,4,3],v=1,eps=0.000001): [RMS_NORM] NMSE = 0.999654538 > 0.000000100 FAIL RMS_NORM(type=f32,ne=[64,5,4,3],v=0,eps=0.000100): [RMS_NORM] NMSE = 0.996574387 > 0.000000100 FAIL RMS_NORM(type=f32,ne=[1920,5,4,3],v=0,eps=0.000100): OK RMS_NORM(type=f32,ne=[64,5,4,3],v=1,eps=0.000100): [RMS_NORM] NMSE = 0.996428322 > 0.000000100 FAIL RMS_NORM(type=f32,ne=[1920,5,4,3],v=1,eps=0.000100): [RMS_NORM] NMSE = 0.996550262 > 0.000000100 FAIL RMS_NORM(type=f32,ne=[64,5,4,3],v=0,eps=0.100000): [RMS_NORM] NMSE = 0.892706561 > 0.000000100 FAIL RMS_NORM(type=f32,ne=[1920,5,4,3],v=0,eps=0.100000): OK RMS_NORM(type=f32,ne=[64,5,4,3],v=1,eps=0.100000): [RMS_NORM] NMSE = 0.893457130 > 0.000000100 FAIL RMS_NORM(type=f32,ne=[1920,5,4,3],v=1,eps=0.100000): [RMS_NORM] NMSE = 0.893824981 > 0.000000100 FAIL

Will try again tomorrow.

Unfortunately had to keep the if condition. block_size here is the work_group_size taken from ggml_sycl_info().max_work_group_sizes[device];

It's fine either way, but just FYI if that's the case then that condition will always evaluate to true, as the max_wg_size on a device will always be greater than the sub-group size.

I wonder why removing the condition breaks all ne[0] = 64 inputs?🤔

Rbiessy · 2025-05-20T14:09:29Z

I looked at some common sizes used with rms_norm and ran a few of them with test-backend-ops perf on LNL iGPU:
Master:

  NORM(type=f32,ne=[4096,1,1,1],v=0,eps=0.100000):                     57344 runs -    18.59 us/run -       32 kB/run -    1.64 GB/s
  NORM(type=f32,ne=[4096,512,1,1],v=0,eps=0.100000):                    6147 runs -   229.37 us/run -    16384 kB/run -   68.12 GB/s
  RMS_NORM(type=f32,ne=[1536,1,1,1],v=0,eps=0.100000):                 81910 runs -    12.92 us/run -       12 kB/run -    0.89 GB/s
  RMS_NORM(type=f32,ne=[1536,512,1,1],v=0,eps=0.100000):               10924 runs -   132.90 us/run -     6144 kB/run -   44.09 GB/s
  RMS_NORM(type=f32,ne=[3072,1,1,1],v=0,eps=0.100000):                 98292 runs -    10.40 us/run -       24 kB/run -    2.20 GB/s
  RMS_NORM(type=f32,ne=[3072,512,1,1],v=0,eps=0.100000):                5462 runs -   191.00 us/run -    12288 kB/run -   61.37 GB/s
  RMS_NORM(type=f32,ne=[4096,1,1,1],v=0,eps=0.100000):                 98292 runs -    10.81 us/run -       32 kB/run -    2.82 GB/s
  RMS_NORM(type=f32,ne=[4096,512,1,1],v=0,eps=0.100000):                6147 runs -   225.46 us/run -    16384 kB/run -   69.32 GB/s

PR:

  NORM(type=f32,ne=[4096,1,1,1],v=0,eps=0.100000):                     57344 runs -    18.88 us/run -       32 kB/run -    1.62 GB/s
  NORM(type=f32,ne=[4096,512,1,1],v=0,eps=0.100000):                    6147 runs -   236.88 us/run -    16384 kB/run -   65.96 GB/s
  RMS_NORM(type=f32,ne=[1536,1,1,1],v=0,eps=0.100000):                 73719 runs -    14.56 us/run -       12 kB/run -    0.79 GB/s
  RMS_NORM(type=f32,ne=[1536,512,1,1],v=0,eps=0.100000):               10924 runs -   140.06 us/run -     6144 kB/run -   41.84 GB/s
  RMS_NORM(type=f32,ne=[3072,1,1,1],v=0,eps=0.100000):                 98292 runs -    10.37 us/run -       24 kB/run -    2.21 GB/s
  RMS_NORM(type=f32,ne=[3072,512,1,1],v=0,eps=0.100000):                5462 runs -   192.81 us/run -    12288 kB/run -   60.79 GB/s
  RMS_NORM(type=f32,ne=[4096,1,1,1],v=0,eps=0.100000):                 98292 runs -    10.91 us/run -       32 kB/run -    2.80 GB/s
  RMS_NORM(type=f32,ne=[4096,512,1,1],v=0,eps=0.100000):                6147 runs -   228.09 us/run -    16384 kB/run -   68.52 GB/s

There is a bit of noise but overall there is some overhead. From what we've seen *_NORM is low enough in the list of kernels sorted by total time spent on the device so it may not be measurable when running whole models. I think it's fine to merge as it is but good to keep in mind.

Rbiessy

Btw @qnixsynapse out of curiosity are you aware of models needing this support or are you working on closing some gaps in the SYCL backend regardless of whether this is used today?

AD2605 · 2025-05-20T14:29:00Z

ggml/src/ggml-sycl/norm.cpp

    const int nwarps = nthreads / WARP_SIZE;
+
+    x   += sample*stride_sample + channel*stride_channel + row*stride_row;


I suspect as we start adding support for non-contiguous layouts, this calculation would become more common, do you think it's worth adding similar to the following in the common.hpp file,

llama.cpp/ggml/src/ggml-sycl/binbcast.cpp

Line 42 in a4090d1

auto calculate_index = [](const std::array<int, 4> & dims, const std::array<int, 4> & strides,

I think this would be useful. There is also similar code for the non-contiguous batched mul_mats.

Added in latest change.

qnixsynapse · 2025-05-20T14:37:04Z

@Rbiessy Ofcourse, closing some gaps in SYCL backend just in case. It is also possible to use different path/kernel for non contiguous inputs if it significantly effects performance.

Rbiessy · 2025-05-22T08:53:45Z

@Rbiessy Ofcourse, closing some gaps in SYCL backend just in case. It is also possible to use different path/kernel for non contiguous inputs if it significantly effects performance.

I think it's fine as it is. There is just one unresolved comment but otherwise looks good to me.

ggml/src/ggml-sycl/common.hpp

ggml-ci

…ernels

ggml-ci

This reverts commit 43be2d6.

ggml-ci

Rbiessy · 2025-05-23T16:32:07Z

ggml/src/ggml-sycl/common.hpp

+/* Helper for Computing the linear offset into an 4-dimensional ggml_tensor given
+per-dimension sizes, strides, and indices */
+template<int N>
+static __dpct_inline__ size_t calculate_offset(const std::array<int, N> & strides, const std::array<int, N> & indices) {


Nit but I don't think it makes sense to mark the function as static in a header, it's just going to be ignored.

Rbiessy · 2025-05-23T16:35:14Z

ggml/src/ggml-sycl/norm.cpp

        item_ct1.barrier(sycl::access::fence_space::local_space);
        mean_var = 0.f;
-        size_t nreduce = nwarps / WARP_SIZE;
+        const size_t nreduce = (nwarps + WARP_SIZE - 1) / WARP_SIZE;


You can use ceil_div here

github-actions bot added ggml changes relating to the ggml tensor library for machine learning SYCL https://en.wikipedia.org/wiki/SYCL - GPU programming language labels May 18, 2025

qnixsynapse closed this May 18, 2025

qnixsynapse deleted the sycl/non_cont_norms branch May 18, 2025 06:42

qnixsynapse restored the sycl/non_cont_norms branch May 18, 2025 06:56

qnixsynapse reopened this May 18, 2025

qnixsynapse marked this pull request as draft May 18, 2025 07:01

qnixsynapse marked this pull request as ready for review May 19, 2025 06:16

qnixsynapse requested review from Alcpz and NeoZhangJianyu May 19, 2025 06:16

Rbiessy reviewed May 19, 2025

View reviewed changes

qnixsynapse force-pushed the sycl/non_cont_norms branch from 4f9b1bc to d5d39b5 Compare May 19, 2025 11:46

AD2605 reviewed May 20, 2025

View reviewed changes

Rbiessy reviewed May 20, 2025

View reviewed changes

AD2605 reviewed May 20, 2025

View reviewed changes

qnixsynapse force-pushed the sycl/non_cont_norms branch from d5d39b5 to 37b9ba6 Compare May 21, 2025 05:14

AD2605 reviewed May 22, 2025

View reviewed changes

ggml/src/ggml-sycl/common.hpp Outdated Show resolved Hide resolved

qnixsynapse added 7 commits May 22, 2025 20:50

SYCL: Add non contiguous input support to norm kernel

08dc327

refactor and add RMS_NORM non contiguous input support

870fe56

ggml-ci

restore subgroup reduction for multi-subgroup thread blocks in norm k…

926cecd

…ernels

Swap grid dims of nsamples and nrows

a18f0ae

ggml-ci

Revert "Swap grid dims of nsamples and nrows"

2676a35

This reverts commit 43be2d6.

restore not required changes

545e47d

ggml-ci

address review comments: change it to more like SYCL

7c9721f

qnixsynapse added 2 commits May 22, 2025 20:50

Use a common function to calculate offset

b6def81

remove wrap around logic for handling broadcasts

ea051b7

qnixsynapse force-pushed the sycl/non_cont_norms branch from 3ad3415 to ea051b7 Compare May 22, 2025 15:21

Rbiessy approved these changes May 23, 2025

View reviewed changes

Rbiessy reviewed May 23, 2025

View reviewed changes

		const int nwarps = nthreads / WARP_SIZE;

		x += samplestride_sample + channelstride_channel + row*stride_row;

SYCL: Add non contiguous support in RMS_NORM and NORM kernels #13611

Are you sure you want to change the base?

SYCL: Add non contiguous support in RMS_NORM and NORM kernels #13611

Uh oh!

Conversation

qnixsynapse commented May 18, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

qnixsynapse commented May 19, 2025

Uh oh!

Rbiessy left a comment

Choose a reason for hiding this comment

Uh oh!

qnixsynapse commented May 19, 2025

Uh oh!

Alcpz commented May 20, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

AD2605 May 20, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Rbiessy commented May 20, 2025

Uh oh!

Rbiessy left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

qnixsynapse commented May 20, 2025

Uh oh!

Rbiessy commented May 22, 2025

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

qnixsynapse commented May 18, 2025 •

edited

Loading

Alcpz commented May 20, 2025 •

edited

Loading

AD2605 May 20, 2025 •

edited

Loading