Attribute based kernel namming

[tdavidcl]

Extension Type

New Feature

Feature Description

Hi,

I was thinking of this the other day, so I'm throwing the idea here for discussions.

Usually we are naming kernels using a class that is passed to the parallel_for

    defaultQueue.submit([&](sycl::handler& cgh) {
      sycl::accessor accA{bufA, cgh, sycl::read_only};
      sycl::accessor accB{bufB, cgh, sycl::read_only};
      sycl::accessor accR{bufR, cgh, sycl::write_only};

      cgh.parallel_for<vector_add_1>(
          sycl::range{dataSize}, [=](sycl::id<1> globalId) {
            accR[globalId] = accA[globalId] + accB[globalId];
          });
    });

Which does work. However, it require the user to always declare the class in advance. Additionally with the advent of free functions the free function commands #922 I'm wondering if we should provide an alternative naming mechanism relying on attributes rather than class. Such as

    defaultQueue.submit([&](sycl::handler& cgh) {
      sycl::accessor accA{bufA, cgh, sycl::read_only};
      sycl::accessor accB{bufB, cgh, sycl::read_only};
      sycl::accessor accR{bufR, cgh, sycl::write_only};

      cgh.parallel_for(
          sycl::range{dataSize}, [=] (sycl::id<1> globalId) [[sycl::named_kernel("vector_add_1")]] {
            accR[globalId] = accA[globalId] + accB[globalId];
          });
    });

or if i borrow a snippet from #922

  khr::launch(myQueue, range<2> { N, M }, [=](id<2> index)  [[sycl::named_kernel("foo")]] {
    size_t i = index[0];
    size_t j = index[1];
    b[i * M + j] = i * 2014 + j * 42;
  });

If not quite sure about the handling of the templated cases though. I see two options:
1 - passing a constexpr const char* to the attribute ?
2 - allow passing extra types to the attribute such as [[sycl::named_kernel("foo", T)]] ?
3 - Why not also accept a class as an argument to allow the usage of classes for naming as done currently ?

One question also is whether we allow having the same name twice or not ? (can be quite hard with linkage i think, but maybe we could allow the same name if it is in the same translation unit)

Is it a good approach ? Maybe ? But making kernel naming simpler for usage with profiling tools would be wonderful as the current class approach tends to generate very long names in some cases or could be more complex for newcomers. We could consider such a attribute as a shortcut for forward declaration + passing the class name if want to keep it simple.

Related Functionality in C++

Not stricly c++ but clang provides the clang::annotate_type attribute which could be use to annotate lambdas:

From (https://github.com/oneapi-src/SYCLomatic/blob/f350c9d684edb1268f5e8fe9eea0666a680060fd/clang/test/SemaCXX/cxx2b-deducing-this.cpp#L318) documentation :

    const auto l16 = [=]() [[clang::annotate_type("foo")]] [[clang::annotate_type("bar")]] {
        return x;
    };

Related Functionality in Other Languages

clCreateKernel(program, "my_kernel", ...); maybe ?

Related SYCL Extensions

#922 Since unless I missed it there is no clear way to name a kernel for profiling tools.

[keryell]

If you want to use some language extensions, why not using CUDA at the first place? 😉

Usually in SYCL we prefer to use normal C++ as much as possible, like wrapper classes or functions to express compile-time properties.
I think DPC++ has something like that already, if I remember correctly.

[tdavidcl]

If you want to use some language extensions, why not using CUDA at the first place? 😉

I have a strong preference for SYCL, this much i can say 😛

Usually in SYCL we prefer to use normal C++ as much as possible, like wrapper classes or functions to express compile-time properties. I think DPC++ has something like that already, if I remember correctly.

Clearly that what is highlighted with the current naming mechanism. To be fair it is only a profiling related issue where the supplied name get sometime wrapped in some templates and make it quite inconvenient in profiling tools. Which strongly wants me to have a sledgehammer attribute to force the name of the kernel. Also i did not see a way to pass a kernel name to the free function commands in #922.

[gmlueck]

Also i did not see a way to pass a kernel name to the free function commands in #922.

This was intentional. Naming a kernel via a template parameter on the launch command introduces a lot of weird corner cases that are not well specified. For example, what happens if the same lambda is used in two launch commands that have different kernel names? Or the converse, when you use the same name for two different lambdas?

If you want a unique name for the kernel, it seems better to define it as a named kernel object (functor). The only language reason to define a kernel name is because you want to query the kernel. Since this is somewhat an advanced feature, it seems reasonable to ask people to define kernels as named function objects in this case.