[SYCL]subgroup shuffle bug when sg_size=32

[luoyu-intel]

Describe the bug

I'm debugging the SYCL backend of llama.cpp.
I found some kernel output -nan when built with Debug. The root cause is that

sycl::select_from_group(g, x,
		target_offset < logical_sub_group_size
		? start_index + target_offset
		: id);

this returns nan when local_id is larger than 16.
I can also reproduce this by using sg.shuffle:
tmp += sg.shuffle(x, i); when i > 16, the result of shuffle is nan. But they both work fine with Release build.

To reproduce

1. checkout llama.cpp: https://github.com/ggerganov/llama.cpp
2. built with SYCL debug config
3. run inference.

Environment

Windows 11
Intel Arc770
Intel(R) oneAPI DPC++/C++ Compiler 2024.1.0 (2024.1.0.20240308)
Target: x86_64-pc-windows-msvc
Thread model: posix

Additional context

No response

[KornevNikita]

I probably rushed into adding the confirmed label. The reproducer is too huge, it's impossible to debug something this way.
@luoyu-intel thanks for investigating. Is it possible for you to provide some smaller reproducer? May be a kernel or so.

[KornevNikita]

Ok, I see there is already a ticket for this in the internal system. @luoyu-intel do you mind closing this issue?