KhronosGroup / SYCL-Docs

SYCL Open Source Specification
Other
109 stars 67 forks source link

Can the types used for kernel names be cv-qualified? #568

Open tahonermann opened 2 months ago

tahonermann commented 2 months ago

Section 5.2, "Naming of kernels", of revision 8 of the 2020 SYCL Specification lists the requirements for naming of SYCL kernels. Those requirements do not prohibit cv-qualified types. (Those requirements also don't require a class type, but the definition of "kernel name" in the glossary states that "A kernel name is a class type"; this might be a separate issue). Is the following program intended to be well-formed? The Intel OneAPI compiler currently accepts it. https://godbolt.org/z/e8enWEEnb.

#include <sycl/sycl.hpp>
struct Kernel;
int main() {
  sycl::queue q;
  q.submit(
    [=](sycl::handler &cgh) {
      cgh.single_task<Kernel>([]() {});
      cgh.single_task<const Kernel>([]() {});
      cgh.single_task<volatile Kernel>([]() {});
      cgh.single_task<const volatile Kernel>([]() {});
    }
  );
}
nliber commented 2 months ago

Good question!

It doesn't even have to be an existing class type (you could put struct Kernel2 without any forward declaration or definition in the angle brackets and it would work, because doing so is a legitimate but useless construct in standard C++).

IIRC, Codeplay used to use the name to generate tables.

Anyway, I'd be fine with either banning it or allowing it, but we should definitely clarify it after getting more input from other implementors.

gmlueck commented 2 months ago

It would be nice if we could simplify the spec wording about kernel names, with the only requirement that it be unique (across all kernels). This would make your example above a correct program.

Another interesting question is a case like this:

struct MyKernel {/*...*/};

struct Name1;
struct Name2;

void foo(sycl::queue q) {
  q.submit([=](sycl::handler &cgh) {
    MyKernel k;
    cgh.single_task<Name1>(k);
  });
  q.submit([=](sycl::handler &cgh) {
    MyKernel k;
    cgh.single_task<Name2>(k);
  });
}

This is the opposite problem. Here there are two invocations of the same kernel MyKernel. However, each invocation gives the kernel a different name. Is that allowed?