`sycl::get_kernel_id<KernelName>()` fails when `KernelName` is only named in a different translation unit
- Dominant language
- LLVM
- Stars
- 1.5k
- Forks
- 854
- Avg merge
- 3d 17h
- Merged PRs (30d)
- 137
Description
## Context
Found while adding a SYCL-CTS conformance test for KhronosGroup/SYCL-CTS#1212, which specifically calls for testing this exact scenario (`get_kernel_id()` for a kernel defined only in a different translation unit).
## Summary
[SYCL 2020 section 4.11.6, "Obtaining a kernel identifier"](
https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#_obtaining_a_kernel_identifier):
> *Preconditions:* The template parameter `KernelName` must be the type kernel name of a kernel that is defined in the SYCL application. [...]
> Applications which call `get_kernel_id()` for a `KernelName` that is not defined are ill formed, and the implementation must issue a diagnostic in this case.
The kernel does not need to be defined (i.e. named in a kernel invocation function such as `handler::single_task()`) in the *same* translation unit as the call to `get_kernel_id()` - only somewhere in the linked application. With DPC++ (`icpx`, Intel oneAPI DPC++/C++ Compiler 2026.1.1), `get_kernel_id()` throws at runtime with `No kernel found with the specified name` when `KernelName` is named only in another translation unit than the one calling `get_kernel_id()`.
Note that cross-TU kernel invocation itself works fine - only the `get_kernel_id()` lookup fails.
There is already a `FIXME` acknowledging incomplete handling of this case in `sycl/include/sycl/kernel_bundle.hpp`:
```cpp
template kernel_id get_kernel_id() {
// FIXME: This must fail at link-time if KernelName not in any available
// translation units.
return detail::get_kernel_id_impl(
detail::CompileTimeKernelInfo.Name);
}
```
## Reproducer
See the [repro.zip](https://github.com/user-attachments/files/30617307/repro.zip) for the reproducer.
## Isolating the cause
Three control experiments, all with the same compiler/flags:
1. `get_kernel_id()` called for a kernel invoked **in the same TU** -> succeeds, returns correct `kernel_id`, kernel runs and produces the expected result.
2. Direct invocation of a kernel named only **in another TU** (no `get_kernel_id()` call at all) -> succeeds, kernel runs and produces the expected result (`result == 42`).
3. `get_kernel_id()` called for a kernel named only **in another TU** (the reproducer above) -> fails with `No kernel found with the specified name`.
This isolates the bug specifically to `get_kernel_id()`'s runtime lookup (`detail::get_kernel_id_impl`) not finding kernels that are registered by a different translation unit than the caller - cross-TU device code linking and kernel invocation otherwise work correctly.
## Environment
- Intel(R) oneAPI DPC++/C++ Compiler 2026.1.1 (2026.1.1.20260724)
- Devices tested: `opencl:cpu` (12th Gen Intel(R) Core(TM) i5-1235U),
`level_zero:gpu` (Intel(R) Iris(R) Xe Graphics)
- Debian GNU/Linux 13.5 x86_64
Contributor guide
Research direction
Start with sycl/include/sycl/kernel_bundle.hpp and trace get_kernel_id() into detail::get_kernel_id_impl. Reproduce the failure from repro.zip, then compare the same-TU and cross-TU control experiments. Done means get_kernel_id() finds a kernel registered in another translation unit without regressing direct cross-TU invocation.
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
- Quiet
- Clarity
- Clearly specified
- Newbie friendliness
- 48/100