Can multiple kernel name types be associated with the same kernel object type?

[tahonermann]

Specification Version

SYCL 2020 (Revision 10)

Section Number(s)

https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:naming.kernels

Issue Description

Discussion in #568 raised a concern regarding whether multiple kernel name types can be associated with the same kernel object type. This concern is tangential to the primary concern for which that issue was reported, so has been moved to this new issue.

See the code example attached to this issue. The initial concern raised in #568 is this:

I don't think code like this should be allowed. MyKernel already has a name, and I can't think of any reason why somebody would want to assign two different names to the same kernel. Assigning two names to the same kernel means that in practice it will have to be compiled twice, even though the body of the kernel will be identical.

... assigning two different names to the same lambda seems broken to me, just like the named function object case.

Further discussion indicated that the suggested duplicate compilation is limited to device entry points (if applicable to the implementation) and SYCL invocation function template instantiations.

The kernel name is what defines the kernel, not the kernel object type. Greg's example requires two distinct device entry points, but I don't think it necessitates multiple compilations of the kernel object type or the functions it calls. Entry points can share common code. Template instantiations would of course be performed for each invocation of the SYCL invocation function with a distinct kernel name type.

Discussion continued with an emphasis on design:

I meant that the design seems broken: it's definitely possible to implement something that allows a user to associate multiple arbitrary type names with the same function object, I just don't understand why that's useful to anybody. It makes the implementation more complicated, means we have to specify the behavior of more weird corner cases, and potentially leads to more confusion for developers.

The code example attached to this issue attempts to demonstrate how a single kernel object type might reasonably used for multiple purposes.

Building off @gmlueck's example: Are get_kernel_id<Name1>() and get_kernel_id<Name2>() required or allowed to return the same kernel_id? Can both Name1 and Name2 be declared in the same kernel bundle? If backend interop is used to extract the underlying OpenCL/Level Zero/CUDA representation of the kernels for Name1 and Name2, can they be the same?

These are questions that should have clear answers in the specification. A useful outcome of this issue discussion might include updating the specification to more clearly specify the correspondence between a kernel ID, a kernel name, and a kernel.

The SYCL specification is not particularly precise in its use of the term "kernel"; see #603.

Code Example (Optional)

struct Kernel {
  bool DoTheOtherThing;
  void operator()() const {
    if (! DoTheOtherThing) { ... }
    else { ... }
  }
};

struct DoTheThing;
struct DoTheOtherThing;

void foo(sycl::queue q) {
  q.submit([=](sycl::handler &cgh) {
    Kernel k { false };
    cgh.single_task<DoTheThing>(k);
  }).wait();
  q.submit([=](sycl::handler &cgh) {
    Kernel k { true };
    cgh.single_task<DoTheOtherThing>(k);
  }).wait();
}

[tahonermann]

Another code example to consider. This is similar to the one from the issue description, but uses a lambda expression instead of a function object type with behavior differentiated via a specialization constant.

constexpr sycl::specialization_id<bool> DoMore;

struct DoTheThing;
struct DoTheOtherThing;

void foo(sycl::queue q, bool b) {
  auto l = [] (sycl::kernel_handler kh) {
    ...
    if (kh.get_specialization_constant<DoMore>()) {
      ...
    }
  };
  if (b) {
    q.submit([=](sycl::handler &cgh) {
      cgh.set_specialization_constant<DoMore>(true);
      cgh.single_task<DoTheThing>(l);
    }).wait();
  } else {
    q.submit([=](sycl::handler &cgh) {
      cgh.set_specialization_constant<DoMore>(false);
      cgh.single_task<DoTheOtherThing>(l);
    }).wait();
  }
}

I reserve my right to not defend the above as an example of good practice! :)

[gmlueck]

I think the presence (or lack) of specialization constants in the kernel is irrelevant. It's normal to have a single name represent a kernel that could be specialized in many different ways:

constexpr sycl::specialization_id<int> variation;

struct MyKernel;

void foo(sycl::queue q, int val) {
  auto l = [] (sycl::kernel_handler kh) {
    ...
    switch (kh.get_specialization_constant<variation>()) {
      ...
    }
  };
  q.submit([=](sycl::handler &cgh) {
    cgh.set_specialization_constant<variation>(val);
    cgh.single_task<MyKernel>(l);
  }).wait();
}

Here the kernel MyKernel can be specialized in any number of ways, but there is still a single name that represents it. There is no expectation that every specialization of the kernel must have a unique name.

[tahonermann]

Here the kernel MyKernel can be specialized in any number of ways, but there is still a single name that represents it. There is no expectation that every specialization of the kernel must have a unique name.

Agreed. The goal with my example was to demonstrate that there might also be reasonable cases where the same kernel object type is used with distinct SYCL kernel invocation functions. I had previously been under the impression that a unique kernel name type was needed for each such invocation, but I see now that DPC++ doesn't require that (or at least doesn't diagnose it). https://godbolt.org/z/1ErEza71x.

What about distinct SYCL kernel invocations for the same kernel object type in different translation units? I don't have a strong sense of how to apply linkage to kernel names and kernel entry points. I guess:

• If the kernel name type has internal linkage, than a synthesized kernel entry point also has internal linkage.
• If the kernel name type has external linkage, than a synthesized kernel entry point has weak/comdat linkage; like an inline function.

[gmlueck]

What about distinct SYCL kernel invocations for the same kernel object type in different translation units? I don't have a strong sense of how to apply linkage to kernel names and kernel entry points.

Kernel names are types. I don't think types in C++ have "linkage", do they? I've been thinking that kernel type-names are "global" in the sense that the same name can be used in different translation units to refer to the same kernel. For example:

// TU1
struct MyKernel;

void foo(sycl::queue q) {
  q.submit([=](sycl::handler &cgh) {
    cgh.set_specialization_constant<variation>(val);
    cgh.single_task<MyKernel>([=]{});
  }).wait();
}

// TU2
struct MyKernel;

void bar() {
  auto id = sycl::get_kernel_id<MyKernel>();
}

Here, I would expect that the call to get_kernel_id<MyKernel>() in TU2 to return the ID of the kernel defined in TU1. Thus, the name MyKernel is "global" across both TUs. As a corollary, I think it would be invalid to define two kernels with the same type-name even if they are defined in different TUs.

Anonymous namespaces present an interesting question. What if the same type-name is defined in two different TUs, each in an anonymous namespace? Logically, you would think that these would be different kernel names, so they wouldn't conflict. However, I'm not sure it would be easy for compilers to implement it this way.

[tahonermann]

Kernel names are types. I don't think types in C++ have "linkage", do they? I've been thinking that kernel type-names are "global" in the sense that the same name can be used in different translation units to refer to the same kernel. For example:

All declared names have linkage ([basic.link]). A type that has external linkage is the same type in every translation unit (and is thus subject to ODR violations).

Here, I would expect that the call to get_kernel_id<MyKernel>() in TU2 to return the ID of the kernel defined in TU1. Thus, the name MyKernel is "global" across both TUs. As a corollary, I think it would be invalid to define two kernels with the same type-name even if they are defined in different TUs.

Agreed. Your example is valid since MyKernel refers to the same type in each translation unit. And yes, associating the same type with different kernels in different TUs should be invalid and considered similar to an ODR-violation. Diagnosing such cases has some challenges since it presumably is valid to use the same kernel name type for invocations of the same kernel in multiple translation units (particularly within inline functions). For a typical tool chain, it is up to the linker to diagnose the violation based on different definitions of common symbols; e.g., different definitions for the device entry point that DPC++ generates.

Anonymous namespaces present an interesting question. What if the same type-name is defined in two different TUs, each in an anonymous namespace? Logically, you would think that these would be different kernel names, so they wouldn't conflict. However, I'm not sure it would be easy for compilers to implement it this way.

Such types are distinct in each TU, so they don't conflict. In common tool chains, symbols that depend on such types are given internal visibility and can't be referenced from other TUs. I suspect notions of internal visibility aren't particularly applicable for SYCL since kernels are exposed in kernel bundles that presumably don't preserve TU associations; at least not in a way that allows sycl::get_kernel_id<MyKernel>() to return different kernel IDs for calls in different translation units when MyKernel does not have external linkage. Differentiating such kernels presumably requires that some unique identifier is generated for each and available at compile-time for association with the TU-local type. EDG supports a module ID (translation unit ID) that can be mangled into a symbol name. If I remember right, it is a hash of file name, file time stamp, and maybe some other data. I'm not sure about other compilers. I'm not sure it is reasonable to require implementation to handle these situations; some investigation into what existing implementations do is warranted.

[tahonermann]

#901 has been filed to follow up on the internal linkage/visibility concerns.

[tahonermann]

@AlexeySachkov raised the following example in an offline discussion.

#include <sycl/sycl.hpp>                            
                                                    
class Functor {                                     
public:                                             
  void operator()(sycl::id<1>) const {}             
  void operator()(sycl::id<2>) const {}             
};                                                  
                                                    
int main() {                                        
  sycl::queue q;                                    
                                                    
  q.submit([&](sycl::handler &cgh) {                
    Functor f;                                      
    cgh.parallel_for(sycl::range<1>(1024), f);      
  });                                               
                                                    
  q.submit([&](sycl::handler &cgh) {                
    Functor f;                                      
    cgh.parallel_for(sycl::range<2>(1024, 1024), f);
  });                                               
}

DPC++ rejects this program as defining two distinct kernels with the same name (DPC++ uses the Itanium ABI typeinfo mangled name for the kernel name type to name a kernel). https://godbolt.org/z/zGvfzoq4n.

<source>:3:7: error: definition with same mangled name '_ZTS7Functor' as another definition
    3 | class Functor {                                     
      |       ^
<source>:3:7: note: previous definition is here

DPC++ accepts the program if one or both of the kernel invocation calls are changed to supply an explicit (and unique) kernel name type.

I don't think the SYCL specification has wording to address this situation.

I think it is clear that the above program must be specified as ill-formed if the same kernel name type (whether explicitly or implicitly provided) is used for each kernel invocation function call. A kernel is defined by the set of all functions reachable from the initial call to a specific call operator of a specific kernel object type. The above program therefore defines two distinct kernels since the set of all reachable functions may differ. While it is reasonable to allow a single kernel to be associated with multiple kernel names, it is not reasonable to allow a single kernel name to be associated with multiple kernels (since calls to sycl::get_kernel_id<>() become ambiguous).

However, the above program can be rendered well-formed if implicit kernel name type selection is required to consider the call operator selected by overload resolution. For example, instead of using the kernel object type as the implicit kernel name type, an implementation could use a class template parameterized by the pointer to member function type of the selected call operator (or just the pointer to member type, but we already have direction via #855 to disallow such types as kernel name types). However, this would complicate use of implicit kernel name types with kernel query functions like sycl::get_kernel_id().

Note that support for kernel object types that declare call operator overload sets effectively necessitates the ability to associate more than one kernel name type with a single kernel object type.

Based on the above, I think the DPC++ behavior is reasonable (though the diagnostic could be improved). To be clear, the reasonableness of its behavior extends to 1) the rejection of the program when the same kernel name type (whether explicitly or implicitly provided) is used for both kernel invocation calls, and 2) acceptance of the program when one or both of the kernel invocation calls specifies a (unique) explicit kernel name type.

At a minimum, the SYCL specification should be updated to clarify the intended behavior.

[Pennycook]

I used to think that the ability to define overloads of operator() within a single function object was a neat way to define flexible kernels, and I suspect I've included that claim in at least one SYCL training presentation. But all of our recent discussions about kernel names have changed my mind; I'm now of the opinion that there really needs to be a 1-to-1 mapping between an obvious user-visible type name and a kernel ID in order for these things to make sense.

Clarifying in the specification that there can only be one such operator() in a kernel function object would force the latest example to be rewritten as:

template <int Dimensions>
class Functor {
public:
  void operator()(sycl::id<Dimensions>) const {}
};

...which I think is actually a lot clearer. Functor<1> and Functor<2> are much more user-friendly names than pointers to member functions would be.

---

This is a little bit of an aside, so feel free to ignore this or tell me to open a separate issue. But one problem I expect us to face if we make Alexey's example ill-formed is that a slightly modified example compiles without any problems:

#include <sycl/sycl.hpp>

class Functor {
public:
  void operator()(sycl::id<1>) const {}
  void operator()(sycl::id<2>) const {} // This is never called, so implementations will ignore it
};

int main() {
  sycl::queue q;

  q.submit([&](sycl::handler &cgh) {
    Functor f;
    cgh.parallel_for(sycl::range<1>(1024), f);
  });                                        
}

...which might lead to surprises. Basically, there's no way for the writer of Functor to check they've written a valid kernel function object in isolation. Programs that use only one operator will be fine, and only programs that use both will fail. As a developer I would find this really weird.

I'm beginning to think that it would be really nice to have some sort of decoration (i.e., a [[sycl::kernel]] attribute, or a class to derive from) that would force a SYCL implementation to compile a function object as a kernel without first seeing a call to a kernel launch function. That would allow implementations to diagnose that a class like Functor has too many operator() overloads.

[PeterTh]

We talked about this in today's WG meeting, and I'll do my best to provide a summary of what was discussed.

The core issue we discussed:
Specifically when referring to kernel functors by the class name of the functor, there is an ambiguity when there is more than one call operator defined (and used, though that doesn't really matter for the overall issue of ambiguity).

Potential solutions:

1. Generally disallow functors which provide multiple call operators from being used as kernels.

• Advantages: clear semantics; relatively easy to diagnose in the front-end (especially with a decoration as pointed out above; but at the latest even without it as soon as any kernel call is made)
• Disadvantages: prevents some types of generic code. Even constructs that don't look all that generic at first glance (e.g. a generic lambda introduced by an auto parameter) could result in this rule being violated, even if the user never wants to refer to the kernel by name.

2. Clarify that only functors with a single call operator can be referred to as kernels via just the class name of that functor; and that the implicit kernel name an implementation generates in cases where there is more than one operator has to encode that operator's signature.

• Advantages: arbitrary generic programs where users do not refer to kernels by names work; existing use of (single-operator) functors referred to by class name also continue to work; users have the flexibility to refer to specific operators of multi-operator functors by providing a manual name at the call site.
• Disadvantages: more complex semantics and implementation; perhaps less clear error messages

Note that all of this is not necessarily actually directly answering the question posed in this issue -- regardless of the solution chosen, there's an independent question of whether a user is allowed to invoke the same kernel object with two different manually specified names.

[tahonermann]

In consideration of potential solution 2 from the previous comment, I spent some time looking into how an implementation could generate a unique type that incorporates both the kernel object type and the signature of the call operator for a given SYCL kernel invocation. Since SYCL is intended to be implementable as a library, obtaining the call operator signature should be implementable using standard C++. The set of arguments to be passed to the call operator is known for each SYCL kernel invocation function, so obtaining the signature of the call operator entails querying the type of the function selected by overload resolution. Unfortunately, C++ does not currently provide a way to deduce the type of a function selected by overload resolution. The following proposal would make such deduction possible, but a solution won't be available until C++29 at the earliest.

• P2825: Overload resolution hook: declcall( unevaluated-call-expression )

In the meantime, an implementation could form a type based on the call signature rather than the type of the called function. However, since conversions are not considered, doing so may result in multiple distinct kernel name types being associated with the same kernel (where a kernel is defined as the invoked call operator and all device functions reachable from it). For example:

#include <sycl/sycl.hpp>

struct convertible_from_item {
  convertible_from_item(sycl::item<1>) {}
  convertible_from_item(sycl::item<2>) {}
  convertible_from_item(sycl::item<3>) {}
};

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

void f(sycl::handler &h, kernel &k) {
  h.parallel_for({ 1 }, k);       // Call operator argument is sycl::item<1>.
  h.parallel_for({ 2, 2 }, k);    // Call operator argument is sycl::item<2>.
  h.parallel_for({ 3, 3, 3 }, k); // Call operator argument is sycl::item<3>.
}

Note that SYCL 2020 section 4.9.4.2.2, "parallel_for invoke", explicitly endorses use of such conversions:

For the simplest case, users need only provide the global range (the total number of work-items in the index space) via a SYCL range parameter. In this case the function object that represents the SYCL kernel function must take one of: 1) a single SYCL item parameter, 2) a single generic parameter (template parameter or auto) that will be treated as an item parameter, 3) any other type implicitly converted from SYCL item, representing the currently executing work-item within the range specified by the range parameter.

DPC++ currently rejects the above example (https://godbolt.org/z/oqcc7aYhf) because the kernel entry point function synthesized for each parallel_for() call differs but is given the same name (_ZTS6kernel corresponding to the mangled name of kernel implicitly used as the SYCL kernel name type for each call). This suggests one of:

1. The definition of a kernel that I've been using is insufficient (the definition must also include any "setup" code needed to invoke the SYCL kernel; e.g. the code that constructs the convertible_from_item argument from one of the sycl::item class template specializations), or
2. The DPC++ behavior is non-conforming (each synthesized kernel entry point should be given a unique name based on the SYCL kernel name).

Tangent: If parallel_for() were to be called for a sycl::kernel object returned by sycl::get_kernel(sycl::get_kernel_id<kernel>()) for the example above, would the caller be required to call sycl::handler::set_args() for an object of type convertible_from_item or for an object of the appropriate sycl::item specialization?

With regard to potential solution 1 from the previous comment, I don't think that solution is viable. Per the reference to section 4.9.4.2.2 above, a call operator may be a template. I'm not aware of anything that prohibits the following example:

#include <sycl/sycl.hpp>

struct kernel {
  template <typename Item>
  void operator()(Item) const {}
};

void f(sycl::handler &h, kernel &k) {
  h.parallel_for({ 1 }, k);
  h.parallel_for({ 2, 2 }, k);
  h.parallel_for({ 3, 3, 3 }, k);
}

DPC++ also rejects this example (https://godbolt.org/z/EenGeb8cP) for the same reason as the previous example; it treats each of the parallel_for() calls as unique kernel invocations, all of which use the same SYCL kernel name type.

Prohibiting multiple call operators in kernel object types would not suffice to reject the above example unless call operator templates were also prohibited. If call operator templates are not prohibited, then implementing a diagnostic at the point of a kernel invocation becomes more challenging.

[tahonermann]

I don't recall discussing whether a kernel object type should be usable to query a kernel even if an explicit kernel name is provided at the point of the kernel invocation. For example, should the following be allowed? This would enshrine the notion of kernels potentially having multiple names.

#include <sycl/sycl.hpp>

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

void f(sycl::queue &q, kernel &k) {
  q.single_task<struct explicit_kernel_name>(k).wait(); // Kernel invocation with explicit kernel name type.
  sycl::get_kernel_id<kernel>();                        // Ok because kernel is implicitly an additional valid kernel name?
}

If that example were to be allowed, presumably the following example would be invalid due to the attempted use of the same kernel name type for distinct kernels.

#include <sycl/sycl.hpp>

struct kernel1 {
  void operator()() const {}
};
struct kernel2 {
  void operator()() const {}
};

void f(sycl::queue &q, kernel1 &k1, kernel2 &k2) {
  q.single_task<kernel2>(k1).wait();  // Kernel invocation of kernel1 using kernel2 as the kernel name type.
  sycl::get_kernel_id<kernel1>();     // Ok because kernel1 is implicitly an additional valid kernel name equivalent to kernel2?
  q.single_task<kernel1>(k2).wait();  // Ok?
                                      // Or invalid because kernel1 is already associated with a different kernel?
                                      // Or invalid because kernel2 is already associated with a different kernel?
  sycl::get_kernel_id<kernel2>();     // Ok because kernel2 is explicitly associated with the kernel from the first call to single_task()?
                                      // Or invalid because kernel2 is associated with multiple kernels?
}

[gmlueck]

We discussed this example in the WG meeting today:

q.single_task(k1);
q.single_task<kernel2>(k1);
q.single_task(k2);
get_kernel_id<kernel2>();

@tahonermann thinks this should be illegal and diagnosed on the third line because kernel2 is used at that point as a duplicate kernel name.