intel / intel/llvm

`sycl::get_kernel_id<KernelName>()` fails when `KernelName` is only named in a different translation unit

Open
#22,852 1 comment 0 reactions 0 assignees View on GitHub
confirmed
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

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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.