[SYCL] Compiler does not exit with error and report unresolved symbols in SYCL device code
- Dominant language
- LLVM
- Stars
- 40.5k
- Forks
- 18.7k
- PR merge metrics
- PR metrics pending
Description
On e9e42884a3255729a0ee53c26d45cc992c412924, Ubuntu 24.04 with Intel PVC or 25.04 with Intel BMG card. Consider the program with SYCL kernel:
```c++
#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]);
__assert_fail("fake assert", "main.cpp", 0, __func__);
});
});
q.wait();
printf("%d\n", data3[0]);
free(data3, q);
free(data2, q);
free(data1, q);
return 0;
}
```
Building the program will succeed:
```console
$ clang++ -fsycl test.cpp -o a.clang
$ echo $?
0
```
But program will fail to launch SYCL kernel:
```console
$ ./a.clang
terminate called after throwing an instance of 'sycl::__V0::exception'
what(): 3 (OL_ERRC_INVALID_BINARY) zeKernelCreate failed with error 2013265944, ZE_RESULT_ERROR_INVALID_MODULE_UNLINKED
Aborted
```
The reason for that is that `__assert_fail` is undefined external symbol which did not end up reported during the program build. Adding `-Wl,--no-undefined` won't help:
```console
$ clang++ -fsycl test.cpp -Wl,--no-undefined -o a.clang
$ echo $?
0
```
The reason of the failure can be debugged with the following patch:
```
diff --git a/offload/plugins-nextgen/level_zero/src/L0Program.cpp b/offload/plugins-nextgen/level_zero/src/L0Program.cpp
index 7786164ae004..c37d27873b42 100644
--- a/offload/plugins-nextgen/level_zero/src/L0Program.cpp
+++ b/offload/plugins-nextgen/level_zero/src/L0Program.cpp
@@ -96,6 +96,18 @@ Error L0ProgramBuilderTy::addModule(size_t Size, const uint8_t *Image,
ze_result_t RC;
CALL_ZE(RC, zeModuleCreate, getZeContext(), L0Device.getZeDevice(),
&ModuleDesc, &Module, &BuildLog);
+
+ size_t szLog = 0;
+ zeModuleBuildLogGetString(BuildLog, &szLog, nullptr);
+
+ char* strLog = (char*)malloc(szLog);
+ zeModuleBuildLogGetString(BuildLog, &szLog, strLog);
+
+ printf(">>>> RC=%d\n", RC);
+ printf("%s\n", strLog);
+
+ free(strLog);
+
if (BuildLog)
zeModuleBuildLogDestroy(BuildLog);
if (RC != ZE_RESULT_SUCCESS) {
```
The output will be:
```console
$ ./a.clang
>>>> RC=0
error : unresolved external symbol __assert_fail at offset 476 in instructions segment #0 (aka kernel : _ZTS10TestKernel)
error : unresolved external symbol __assert_fail at offset 492 in instructions segment #0 (aka kernel : _ZTS10TestKernel)
terminate called after throwing an instance of 'sycl::__V0::exception'
what(): 3 (OL_ERRC_INVALID_BINARY) zeKernelCreate failed with error 2013265944, ZE_RESULT_ERROR_INVALID_MODULE_UNLINKED
Aborted
```
**Can SYCL linkage be fixed by reporting errors if unresolved symbols present in device code?**
CC: @tahonermann, @YuriPlyakhin
Contributor guide
Research direction
Reproduce the issue with the provided SYCL program and clang++ -fsycl, then inspect offload/plugins-nextgen/level_zero/src/L0Program.cpp around L0ProgramBuilderTy::addModule and zeModuleCreate. Use the Level Zero build log shown in the report to trace how unresolved device symbols are handled. Done means the unresolved __assert_fail is reported as a build error rather than allowing a binary that fails at kernel launch.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- compilers, devtools
- Issue type
- Bug
- Difficulty
- 4/5
- Estimated time
- 3-5 days
- Activity status
- Active
- Clarity
- Mostly clear
- Newbie friendliness
- 50/100