KhronosGroup / KhronosGroup/SYCL-Docs
Are implementations required to provide implicit kernel names for kernel object types that are not forward declarable?
- Dominant language
- JavaScript
- Stars
- 158
- Forks
- 80
- Avg merge
- 7d 6h
- Merged PRs (30d)
- 5
Description
### Specification Version
SYCL 2020 (Revision 9)
### Section Number(s)
https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#subsec:invokingkernels
https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:naming.kernels
https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:feature-sets.reduced
### Issue Description
Consider the following program that passes a kernel object of local class type to a SYCL invocation function without specifying an explicit kernel name type.
```
#include
int main() {
struct kernel {
void operator()() const {}
};
sycl::queue q;
q.submit([](sycl::handler &h){
h.single_task(kernel{});
});
}
```
SYCL 2020 [section 4.9.4.2, "SYCL functions for invoking kernels"](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#subsec:invokingkernels), states:
> Each function takes an optional kernel name template parameter. **The user may optionally provide a [kernel name](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#kernel-name), otherwise an implementation-defined name will be generated for the kernel.**
[Section 5.2, "Naming of kernels"](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:naming.kernels), includes the following in its list of kernel naming requirements:
> - If the kernel is defined as a named function object type, the name can be the typename of the function object as long as it is either declared at namespace scope, or does not conflict with any name in an enclosing namespace scope.
The SYCL specification seems clear that the above program should be accepted, but supporting it poses implementation challenges. The DPC++ oneAPI compiler accepts it in its default full feature set mode (which does not enforce the forward declarability kernel name requirements of section 5.2), but rejects it when operating in its reduced feature set mode (as is required to support use of a standard C++ compiler as a SYCL-unaware host compiler); see https://godbolt.org/z/hdoaKonoG. When invoked with the `-fno-sycl-unnamed-lambda` option, the compiler uses the type of the kernel object as the implicit kernel name and relies on an integration header to provide a forward declaration, but `main()::kernel` is not a forward declarable type.
Generation of an implicit unique kernel name type for an example like the one above is difficult, and perhaps impossible, for an implementation that cannot use a type that isn't forward declarable as a kernel name type. The challenge is that the only semantic information available to uniquely identify the defined kernel is the type of the kernel object. Though a class template is forward declarable, a class template specialization parameterized by a type that isn't forward declarable is not forward declarable. The only implementation strategy I can think of would require clever use of [C++ friend injection](https://stackoverflow.com/questions/79520873/c-friend-injection-how-does-it-work-and-what-are-the-rules) to implement a type counter. Given differences in template instantiation order across C++ compilers, I suspect such an implementation is not viable when a standard C++ compiler is used as a SYCL-unaware host compiler.
This issue has been filed to confirm whether conforming SYCL implementations are required to accept the above program and, if they aren't, to clarify the SYCL specification; perhaps by modifying section 4.9.4.2 to clarify when an implementation is obligated to provide an implicit kernel name type and/or amending [section B.2, "Reduced feature set"](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:feature-sets.reduced), to require the programmer to provide an explicit name when the type of the kernel object is not forward declarable.
### Code Example (Optional)
_No response_
Contributor guide
No contributing guide indexed for this repository
Assessment
This issue has not been assessed yet.