SYCL compilation flag -fsycl-id-queries-fit-in-int does not work as expected for TriuTril kernel.
- Dominant language
- Python
- Stars
- 113
- Forks
- 128
- Avg merge
- 5d 9h
- Merged PRs (30d)
- 112
Description
This issue is created to track unexpected behavior observed when launching sycl kernels on XPU.
## Problem description
During SYCL code build the flag `-fsycl-id-queries-fit-in-int` is the default one: [DOCUMENTATION](https://www.intel.com/content/www/us/en/docs/dpcpp-cpp-compiler/developer-guide-reference/2025-2/fsycl-id-queries-fit-in-int.html)
In the main PyTorch build system there is no such flag, so the default one is used.
It assumes that the ranges passed when launching sycl kernel fit into `uint32` limits. That supposedly allows some optimizations by not using double precision for indexing.
I created simple cpp sycl kernel to test it.
The attempt to launch sycl kernel with range exceeding max `uint32` ends up in thrown Exception.
```cpp
size_t large_range = 5000000000UL; // > uint32 max
q.submit([&](sycl::handler& h) {
h.parallel_for(sycl::range<1>(large_range), [=](sycl::id<1> idx) {
auto global_id = idx[0]; // Throws exception with default flag
// kernel work...
});
});
```
Exception:
```text
Provided range and/or offset does not fit in int. Pass '-fno-sycl-id-queries-fit-in-int' to remove this limit.
```
On the other hand, when launching the TriuTril kernel in PyTorch ([MENTIONED KERNEL](https://github.com/intel/torch-xpu-ops/blob/97b093e30047055d359e6247d9f80ef9eb94edf5/src/ATen/native/xpu/sycl/TriangularOpsKernels.cpp)), there is no Exception, but the results are wrong.
So in normal, simple case, there is Exception thrown when range exceeds the max `uint32`. We need to disable the default flag to make it work by adding `-fno-sycl-id-queries-fit-in-int` to build.
In PyTorch Triangular kernels for XPU, there is no Exception thrown. Kernel executes, but results are wrong. The experiments showed that the wrong results are caused by `get_global_id()` function that returns the ids truncated to `uint32`.
```cpp
IndexType linear_idx = item.get_global_id(0) * elements_per_thread;
// The linear_idx is effectively:
// linear_idx = real_id % MAX_UINT32
```
## Considerations
I believe it would be good to investigate this mismatch in behavior and find the root cause. It is much better to have the Exception rather than silently truncating `uint64` values to `uint32`.
On the other hand, maybe the flag `-fno-sycl-id-queries-fit-in-int` may be added to build system to avoid the truncation in the kernel and allow ranges bigger than max `uint32`. But due to [DOCUMENTATION](https://www.intel.com/content/www/us/en/docs/dpcpp-cpp-compiler/developer-guide-reference/2025-2/fsycl-id-queries-fit-in-int.html) it may introduce performance issues.
The main goal of this issue is to first find why the TriangularOpsKernels doesn't throw Exceptions when range is bigger than max `uint32` value.
## Additional links
There is the discussion about it in one of the PRs: #2864
Contributor guide
Assessment
This issue has not been assessed yet.