Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

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

Open
tahonermann opened this issue Jun 20, 2024 · 2 comments
Open

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

tahonermann opened this issue Jun 20, 2024 · 2 comments

Comments

@tahonermann
Copy link

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
Copy link
Collaborator

nliber commented Jun 20, 2024

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
Copy link
Contributor

gmlueck commented Jun 24, 2024

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?

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

No branches or pull requests

3 participants