intel / intel/torch-xpu-ops

SYCL compilation flag -fsycl-id-queries-fit-in-int does not work as expected for TriuTril kernel.

Open
#2,950 0 comments 0 reactions 1 assignee Claimed by @BBBela View on GitHub
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

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.