llvm / llvm/llvm-project

[SYCL] Compiler does not exit with error and report unresolved symbols in SYCL device code

Open
#221,283 2 comments 0 reactions 0 assignees View on GitHub
confirmed SYCL
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

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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.