Types in the std namespace do not need to be prohibited as kernel names due to forward declaration requirements

[tahonermann]

SYCL 2020 section 5.2, "Naming of kernels" includes the following note:

The requirement that a kernel name be forward declarable makes some types for kernel names illegal, such as anything declared in the std namespace (adding a declaration to namespace std leads to undefined behavior).

It is correct that C++ programs are prohibited from adding declarations in the std namespace ([namespace.std]), but that does not imply that an existing type in the std namespace cannot be used as a kernel name.

#include <sycl/sycl.hpp>
#include <vector>
int main() {
  sycl::queue q;
  q.submit(
    [=](sycl::handler &cgh) {
      cgh.single_task<std::vector<bool>>(
        [=]{}
      );
    }
  );
  q.wait();
}

The requirement for forward declarable types is presumably intended to accommodate implementations like DPC++ that use the integration header technique to make kernel details discovered during device compilation available during host compilation. In that case, the integration header is a detail of the implementation and not subject to the prohibition on adding declarations to the std namespace (an argument can be made that the SYCL library implementation is distinct from the C++ implementation and is therefore not exempt from the std namespace rule, but DPC++ at least does not require that distinction).

To be clear, I don't think there are any benefits to be gained by allowing types declared in the std namespace to be used as kernel names; I just disagree with the explanation in the referenced note and would prefer to see the note replaced with a normative requirement that simply prohibits the use of types (directly or indirectly) declared in the std namespace from being used as kernel names.

[tahonermann]

There is a wider concern here. The current specification appears to prohibit entities from the std namespace from being used as template arguments of a parameterized kernel name. From SYCL 2020 section 5.2, "Naming of kernels":

• 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.
• ...
• If a kernel function relies on template parameters, then those template parameters must be contained by the kernel name. If such a kernel name is specified as a template argument in a kernel invoking interface, then the template parameters on which the kernel depends must be forward declarable at namespace scope.

Should implementations therefore reject the following program because it uses std::byte as a template argument in the kernel name?

#include <sycl/sycl.hpp>

template<typename T>
struct kernel {
  kernel(T t) : t(t) {}
  void operator()() const {}
  T t;
};

int main() {
  sycl::queue q;
  q.submit(
    [=](sycl::handler &h) {
      kernel k(std::byte(42));
      h.single_task<decltype(k)>(k); // decltype(k) is kernel<std::byte>.
    }
  );
  q.wait();
}

[nliber]

IIRC Codeplay used to use the name for tables its compiler generated (but they'd know for sure).

I think I'd rather we go in the direction of forward declared without a definition (although it may be too late for this, and I don't think this is worth breaking code over). After all, it's just supposed to be a unique name. That covers things like h.single_task<class MyType>(...) as well, where MyType has not had a declaration until now.

I don't expect compilers will be required to warn on this.

[tahonermann]

@nliber, I can't speak for Codeplay, but DPC++and the Intel compiler currently implement the integration header approach and thus depend on forward declaration injection; at least for support of third party host compilers.

I think I'd rather we go in the direction of forward declared without a definition

I don't see what that would help. It seems that would require programmers to invent more kernel names than they need to now.

I don't expect compilers will be required to warn on this.

I'm working on implementing warnings for invalid kernel name types in general right now. The current user experience provided by DPC++ and the Intel compiler when invalid kernel name types are used is ... not easy to explain. https://godbolt.org/z/K7Ev948bE.

[gmlueck]

I agree that it would be good to clarify this language about legal kernel type-names. The term "forward declarable" has never seemed very precise to me. Instead, I wonder if we can make a list of legal types, similar to what we did for the Intel extension sycl_ext_oneapi_free_function_kernels.

I agree with @tahonermann's comment above about std::byte. I think people would be surprised if they were unable to use std::byte as a template parameter for a named kernel function. Along this same line, people would be surprised if we disallowed type aliases in the std namespace like std::size_t or std::uint32_t. I wonder if we should just drop the prohibition against allowing types in the std namespace.

[tahonermann]

I wasn't going to bring up the free function kernel proposal, but since @gmlueck now has, well, I'll widen the scope a bit more. Parameters of free function kernels also need to be forward declarable for implementations that use the integration header technique. If entities in the std namespace cannot be forward declared, then that would effectively prohibit use of the standard library in the interface of such functions (with the exception of type aliases; forward declarations can use the target type of the alias and omit any mention of the alias).