Remove implementation-defined generation of a kernel name type

[tahonermann]

Specification Version

SYCL 2020 (Revision 10)

Section Number(s)

https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#subsec:invokingkernels

Issue Description

SYCL 2020 section 4.9.4.2, "SYCL functions for invoking kernels" states:

Each function takes an optional kernel name template parameter. The user may optionally provide a kernel name, otherwise an implementation-defined name will be generated for the kernel.

Discussion in the SYCL WG in the context of #568 and #893 has revealed that programmers expect to be able to use the type of the kernel object as the kernel name type when a kernel name type is not explicitly provided. For example, programmers expect the following example to be well-formed and to behave the same for all implementations.

#include <sycl/sycl.hpp>

struct kernel {
  void operator()() const {}
};

int main() {
  kernel k;
  sycl::queue q;
  q.single_task(k);
  q.wait();
 auto kid1 = sycl::get_kernel_id<kernel>();      // Ok; retrieves the kernel ID for the invocation of k.
 auto kid2 = sycl::get_kernel_id<decltype(k)>(); // Ok; retrieves the kernel ID for the invocation of k.
}

This expectation extends to lambda closure types and it has been reported that kokkos contains code like the following:

int main() {
  auto k = []{};
  sycl::queue q;
  q.single_task(k);
  q.wait();
 auto kid2 = sycl::get_kernel_id<decltype(k)>(); // Ok; retrieves the kernel ID for the invocation of k.
}

This expectation seems reasonable subject to the reduced feature set limitations imposed on use of lambda expressions without an explicit kernel name type and the forward declarable kernel name type concerns raised in #857 and #858. If the forward declarable kernel name type requirements were moved to the reduced feature set as suggested in #859, then there would appear to be no technical reason not to specify that the kernel object type is always used as the implicit kernel name type for full feature implementations.

Code Example (Optional)

No response

[gmlueck]

This expectation extends to lambda closure types and it has been reported that kokkos contains code like the following:

Just pointing out that this relates to #901. In particular, that issue notes that the type of a lambda closure might be the same in two different translation units, even if the body of the closure (i.e. the code of the kernel) is different. Here's the relevant part of #901:

The issue demonstrated above is applicable to (unnamed) lambda kernels as well. If the example is changed such that k1() and k2() instead defer the kernel invocation to a static function with the same name in each translation unit and that invokes an unnamed lamdba (such that the lambda closure type is given the same (nested) name in each translation unit), DPC++ produces the same behavior.