intel / intel/llvm

SYCL kernels fail to execute if built with DPCLANG -fno-gpu-rdc and linked with GCC

Open
#23,097 14 comments 0 reactions 0 assignees View on GitHub
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

Open the contributing 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.