SYCL kernels fail to execute if built with DPCLANG -fno-gpu-rdc and linked with GCC
- Dominant language
- LLVM
- Stars
- 1.5k
- Forks
- 854
- Avg merge
- 3d 17h
- Merged PRs (30d)
- 137
Description
On https://github.com/intel/llvm/commit/6c46af6bec04bd2d4a8be941197252e1d333f8b7, Ubuntu 24.04, PVC.
I was attempting to build Pytorch with the `-fno-gpu-rdc` option and observed that built GPU kernels fail to be executed. Below is a simple reproducer.
Consider SYCL kernel (it just allocates 2 integers, sums them and writes result in the 3d integer):
```
#include
#include
class TestKernel;
int main() {
sycl::queue q;
const size_t dataSize = 1;
uint8_t *data1 = sycl::malloc_shared(1, q);
uint8_t *data2 = sycl::malloc_shared(1, q);
uint8_t *data3 = sycl::malloc_shared(1, q);
data1[0] = 120;
data2[0] = 0;
data3[0] = 0;
q.submit([&](sycl::handler &cgh) {
cgh.parallel_for(
sycl::range<1>(dataSize),
[=](sycl::id<1> idx) { data3[idx] = std::plus<>{}(data1[idx], data2[idx]); });
});
q.wait();
printf("%d\n", data3[0]);
free(data3, q);
free(data2, q);
free(data1, q);
return 0;
}
```
It can be built and executed as follows:
```
$ dpclang++ -fsycl -fPIE test.cpp -o a.dpclang
$ ./a.dpclang
120
```
Now the same kernel will fail if dpclang++ compiler used just for the compilation stage (with the `-fno-gpu-rdc`) and linkage done with the GCC:
```
$ dpclang++ -fsycl -fno-gpu-rdc -fPIE -c test.cpp -o test.o
$ g++ test.o $(pkg-config --libs sycl-dpcpp-7) -o a.gcc
$ ./a.gcc
a.gcc: /home/dvrogozh/git/intel-llvm/sycl/source/detail/program_manager/program_manager.cpp:1620: sycl::_V1::detail::DeviceKernelInfo& sycl::_V1::detail::ProgramManager::getDeviceKernelInfo(const sycl::_V1::detail::compile_time_kernel_info_v1::CompileTimeKernelInfoTy&): Assertion `It != m_DeviceKernelInfoMap.end()' failed.
Aborted (core dumped)
```
With ICPX oneAPI 2026.1 the same causes segfault in pretty the same place where intel/llvm fails. That's expected as ICPX is built without assertions.
Per my understanding of `-fno-gpu-rdc` above trick should have worked. Please, fix.
Note that there will be no issue if dpclang++ is also used for linkage:
```
$ dpclang++ -fsycl -fno-gpu-rdc -fPIE -c test.cpp -o test.o
$ dpclang++ -fsycl test.o -o a.dpclang2
$ ./a.dpclang2
120
```
CC: @YuriPlyakhin, @KseniyaTikhomirova, @tahonermann
Contributor guide
Research direction
Reproduce the failure using the dpclang++ compilation followed by GCC linkage shown in the issue. Start at sycl/source/detail/program_manager/program_manager.cpp:1620 and compare the failing GCC-linked path with the working dpclang++-linked command; done means the kernel executes successfully when compiled with -fno-gpu-rdc and linked with GCC.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- compilers
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Active
- Clarity
- Mostly clear
- Newbie friendliness
- 48/100