Optimisation bug using sycl::select with sycl::short3

[emtKazys]

Describe the bug

I'm getting zero values when using sycl::select() with sycl::short3 and converting the result to a sycl::id<3>. The following function call will unexpectedly return zero values:

sycl::id<3> wrap_relative_index(sycl::short3 idx, const sycl::short3& dimensions) noexcept
{
  // Wrap underflow indices
  idx = sycl::select(idx, idx + dimensions, idx < 0);
  // Wrap overflow indices
  idx = sycl::select(idx, idx - dimensions, idx >= dimensions);
  return {static_cast<size_t>(idx[0]), static_cast<size_t>(idx[1]), static_cast<size_t>(idx[2])};
}

I can't really distil down a simple explanation, but I do have a single source reproducible case attached along with a Dockerfile to setup a repeatable environment and build the source.

The issue can be worked around by either using sycl::int3 instead of sycl::short3 or by replacing the use of sycl::select() in the example source file (build with -DUSE_INT3 or -DNO_SELECT respectively.

To reproduce

Reproduce with docker file:

1. Install docker and nvidia-container-toolkit
2. Download source file and docker file (attached)
3. Build container docker build -f Dockerfile.txt -t short3_select_bug .
4. Run the container docker run --rm -it --gpus all short3_select_bug bash
5. Run the example: /root/bin/select_short3_bug
6. Run the workarounds:

• /root/bin/select_short3_bug_o0
• /root/bin/select_short3_bug_int3
• /root/bin/select_short3_bug_no_select

select_short3_bug.cpp
Dockerfile.txt

Manual build:

1. Download sample source (above)
2. Open shell and navigate to the source location.
3. source /opt/intel/oneapi/setvars.sh
4. Create output directory: mkdir bin
5. Build: icpx select_short3_bug.cpp -fsycl -sycl-std=2020 -std=c++20 -fsycl-targets=nvptx64-nvidia-cuda -O3 -o bin/select_short3_bug (see also the Dockerfile for workaround variations)
6. Run with bin/select_short3_bug

Expected output:

(0,0,0) => (2,2,2)      expect: (2,2,2) ok? true
(-1,-1,-1) => (1,1,1)   expect: (1,1,1) ok? true
(4,4,4) => (2,2,2)      expect: (2,2,2) ok? true
(5,5,5) => (3,3,3)      expect: (3,3,3) ok? true
(4,4,-2) => (2,2,0)     expect: (2,2,0) ok? true
(-3,-3,-4) => (3,3,2)   expect: (3,3,2) ok? true
(-3,3,1) => (3,1,3)     expect: (3,1,3) ok? true
(0,1,4) => (2,3,2)      expect: (2,3,2) ok? true
(4,0,1) => (2,2,3)      expect: (2,2,3) ok? true
(4,2,-1) => (2,0,1)     expect: (2,0,1) ok? true
(1,-3,2) => (3,3,0)     expect: (3,3,0) ok? true
(4,4,-4) => (2,2,2)     expect: (2,2,2) ok? true
(-4,-2,-4) => (2,0,2)   expect: (2,0,2) ok? true
(-1,-2,3) => (1,0,1)    expect: (1,0,1) ok? true
(1,1,3) => (3,3,1)      expect: (3,3,1) ok? true
(4,3,1) => (2,1,3)      expect: (2,1,3) ok? true

Actual output:

(0,0,0) => (0,0,0)      expect: (2,2,2) ok? false
(-1,-1,-1) => (0,0,0)   expect: (1,1,1) ok? false
(4,4,4) => (0,0,0)      expect: (2,2,2) ok? false
(5,5,5) => (0,0,0)      expect: (3,3,3) ok? false
(4,4,-2) => (0,0,0)     expect: (2,2,0) ok? false
(-3,-3,-4) => (0,0,0)   expect: (3,3,2) ok? false
(-3,3,1) => (0,0,0)     expect: (3,1,3) ok? false
(0,1,4) => (0,0,0)      expect: (2,3,2) ok? false
(4,0,1) => (0,0,0)      expect: (2,2,3) ok? false
(4,2,-1) => (0,0,0)     expect: (2,0,1) ok? false
(1,-3,2) => (0,0,0)     expect: (3,3,0) ok? false
(4,4,-4) => (0,0,0)     expect: (2,2,2) ok? false
(-4,-2,-4) => (0,0,0)   expect: (2,0,2) ok? false
(-1,-2,3) => (0,0,0)    expect: (1,0,1) ok? false
(1,1,3) => (0,0,0)      expect: (3,3,1) ok? false
(4,3,1) => (0,0,0)      expect: (2,1,3) ok? false

Environment

• OS: Linux 24.04
• Target device and vendor: Nvidia GPU
• DPC++ version: 2025.2 also tested at tag v6.1.0
• Dependencies version: [cuda:gpu][cuda:0] NVIDIA CUDA BACKEND, NVIDIA RTX A3000 Laptop GPU 8.6 [CUDA 12.9]

Additional context

No response