`sycl::get_kernel_id<KernelName>()` cannot be implemented in strictly conforming C++

[tahonermann]

Specification Version

SYCL 2020 (Revision 11)

Section Number(s)

SYCL 2020 section 4.11.6, "Obtaining a kernel identifier"
SYCL 2020 section 3.12.3, "Library-only implementation"

Bug Description

SYCL 2020 section 4.11.6, "Obtaining a kernel identifier" states:

Preconditions: The template parameter KernelName must be the type kernel name of a kernel that is defined in the SYCL application. [...] Applications which call get_kernel_id() for a KernelName that is not defined are ill formed, and the implementation must issue a diagnostic in this case.

SYCL 2020 section 3.12.3, "Library-only implementation" states:

It is also possible to implement SYCL purely as a library, using an off-the-shelf host compiler with no special support for SYCL. In such an implementation, each kernel may run on the host system.

As far as I've been able to tell, C++17 and later C++ standards do not provide features sufficient to implement these requirements when the above wording is interpreted as permitting a call to sycl::get_kernel_id<KernelName>() to succeed for a kernel that is defined in a translation unit other than the one with the call while also requiring a diagnostic if the requested kernel is not defined in any translation unit.

The closest C++17 solution I've been able to find is prototyped at https://godbolt.org/z/oW6srrYvj. The prototype exercises the following:

• Requesting a kernel ID for a kernel defined in another translation unit that, at run-time, is submitted after the request for the kernel ID (k1.cpp, k2.cpp).
• Requesting a kernel ID for a kernel defined later in the translation unit that, at run-time, is submitted after the request for the kernel ID (k3.cpp).

Validation that a diagnostic is issued if a requested kernel is not defined in any translation unit can be demonstrated by commenting out the call to h.single_task<KN>([]{}) in k2.cpp.

The prototype relies on class template instantiations and hidden friend function definitions to implement kernel registration across translation units. The basic idea is that, when a SYCL kernel invocation function is called, a class template (kernel_registrar) containing a hidden friend function definition (get_kernel_info_helper()) is instantiated and the friend function is ODR-used. When sycl::get_kernel_id<KN>() is called, a different class template (kernel_registry) containing a friend declaration without a definition for the same function is instantiated and the friend function is ODR-used. Note that hidden friend function definitions are implicitly inline functions. If no definition of the friend function is present in the current translation unit (because no kernel was registered with the requested kernel name type in that translation unit), then a linker error will occur unless a definition is provided in some other translation unit (e.g., one that registered a kernel with that kernel name type). If such a definition is available, then the call to the friend function returns the kernel info object for the registered kernel.

As far as I can tell, this approach conforms with C++17 in all but one respect. C++17 [basic.def.odr]p4 (<basic.def.odr>p13 in the current C++ working draft) states:

... An inline function or variable shall be defined in every translation unit in which it is odr-used outside of a discarded statement.

This requirement is intentionally violated by the prototype in order to provoke a linker error due to a missing symbol when sycl::get_kernel_id<KN>() is called for a type KN that does not correspond to a kernel in any translation unit. This requirement grants implementations permission to not emit a definition for an inline function that is ODR-used in a translation unit because the implementation can rely on a definition being available for any other translation unit that ODR-uses the function. The prototype depends on the translation unit in which a kernel is registered providing a definition for use by translations units that request a kernel ID. GCC and Clang do take advantage of the opportunity to not emit a definition of the inline friend function when optimization is enabled. The prototype uses the GNU __attribute__((used)) attribute to disable that optimization. MSVC does not appear to take advantage of this opportunity even when whole program optimization and link time code generation is enabled.

The code compiles cleanly for Clang and MSVC. GCC issues a rather annoying warning about a friend declaration declaring a non-template function; see the documentation for -Wno-non-template-friend here. The prototype runs as expected under all optimization levels and with link time optimization enabled.

There are a few options for how to address this in the SYCL specification:

1. Prove me wrong by finding a way to implement the stated requirements with a strictly conforming C++17 implementation.
2. Leave the SYCL specification as is with an acknowledgement that the stated requirements cannot be met with strictly conforming C++17 code.
3. Change the SYCL specification to remove the requirement for a diagnostic if the requested kernel is not defined in any translation unit. sycl::get_kernel_id<KN>() could then be changed to result in undefined behavior if KN is not the name of a defined kernel in any translation unit (this would permit implementations to (continue to) reject a program at link-time with a diagnostic) or to throw an appropriate exception. Note that the ability to throw an exception would likely require an implementation to build a kernel registry as part of process startup.

[tahonermann]

Intel implemented the approach described in the issue description and was rewarded with a report of a regression for code like the following.

#include <sycl/sycl.hpp>

template<typename T>
struct KN {
  static_assert(sizeof(T));
};

void f(sycl::handler &h) {
  h.single_task<KN<void>>([]{});
}

The type used for the kernel name type in f() is a class template specialization with a definition that is ill-formed if instantiated. Kernel name types may be incomplete types, so a definition for a kernel name type is not required. What is happening is that the prototype implementation uses the kernel name type such that it is an associated class (C++17 [basic.lookup.argdep]p2; an associated entity ( <basic.lookup.argdep>p3) in the current C++ working draft) in an unqualified function call. Argument Dependent Lookup (ADL) implicitly instantiates the definitions of all associated classes in its search for friend function declarations that might contribute to the overload set for the call. The ADL triggered instantiation of the definition of KN<void> fails because sizeof cannot be applied to an incomplete type. The failure is demonstrated at https://godbolt.org/z/GK1eKcPs6 (for Clang; GCC only issues a warning by default).

Fortunately, it is possible to avoid the interaction with ADL and the undesired instantiation of kernel name type definitions by smuggling the kernel name type through the unqualified function call as a member type alias of a templated class that, though dependent, does not have the kernel name type as an associated class.

An updated prototype that uses this technique is available at https://godbolt.org/z/Ehqj1zf98; see the updated comments in sycl.hpp for details. The prototype works as expected for Clang, GCC, and MSVC (tested locally with Visual Studio 2022 since godbolt.org times out during CMake configuration with MSVC).