Revert "[SYCL] Enable referencing a kernel from a different translation unit (#22264)

sycl/include/sycl/detail/kernel_desc.hpp

@@ -188,94 +188,6 @@ template <typename KNT> struct KernelIdentity {
   using type = KNT;
 };
 
-// KernelRegistrar registers a kernel for later lookup and production of a
-// kernel_id via sycl::get_kernel_id<KN>(). Registration is accomplished by
-// instantiating a specialization with the kernel name type and KernelInfo
-// for each kernel and then calling the registerKernelName() static
-// member function to provoke an ODR-use of the getKernelNameHelper()
-// hidden friend function definition. This makes the definition of
-// getKernelNameHelper() for a particular kernel name available for use
-// by arbitrary translation units without requiring registration in each
-// of them.
-template <typename KN, typename KernelInfo> struct KernelRegistrar {
-  // getKernelNameHelper() provides an interface to retrieve a kernel
-  // entry point name corresponding to a kernel name type. Though dependent
-  // on the KN template parameter, this function is not a member of the
-  // KernelRegistrar class; it is a namespace scope function that can be
-  // redeclared in other class template contexts. Since it is only dependent
-  // on KN, note that multiple instantiations of KernelRegistrar with the
-  // same kernel name but different kernel entry point name symbol would
-  // result in duplicate definitions of getKernelNameHelper(). This is
-  // by design.
-  //
-  // Hidden friend function definitions are implicitly inline. C++17
-  // [basic.def.odr]p4 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 in order to provoke a
-  // linker error when a call to sycl::get_kernel_id<KN>() is made for
-  // a type that has not been registered as a kernel name type; see
-  // the comments associated with the KernelRegistry class definition.
-  // The quoted statement allows implementations to not emit definitions
-  // of inline functions even if ODR-used within a translation unit (e.g.,
-  // because all such uses were inlined). In practice, some implementations
-  // do skip emitting such definitions subject to compile-time and link-time
-  // optimization. The GNU 'used' attribute is applied to force GCC and
-  // Clang to emit such definitions. MSVC does not appear to suppress such
-  // definitions, even with whole program analysis (/WP) and link time
-  // code generation (/LTCG).
-  //
-  // Name lookup for getKernelNameHelper() requires argument dependent
-  // lookup (ADL); even from within this class definition. The call to
-  // getKernelNameHelper() in registerKernelName() therefore passes an
-  // instance of this class (dependent on both KN and KernelInfo) to enable
-  // lookup in this class. A conversion operator then enables the argument
-  // to be converted to the KernelIdentity wrapper type (only dependent on KN).
-#if defined(__GNUG__) || defined(__clang__)
-  __attribute__((used))
-#endif
-#if defined(_WIN32)
-  __declspec(dllexport)
-#endif
-  friend const char *
-  getKernelNameHelper(KernelIdentity<KN>) noexcept {
-    return KernelInfo::getName();
-  }
-  operator KernelIdentity<KN>() const noexcept { return KernelIdentity<KN>{}; }
-  static void registerKernelName() noexcept {
-    (void)getKernelNameHelper(KernelRegistrar<KN, KernelInfo>{});
-  }
-};
-
-// KernelRegistry enables retrieval of a kernel entry point name given
-// kernel name type KN. Lookup is performed by instantiating a
-// specialization with the kernel name type and then calling the
-// getKernelName() static member function. getKernelName() then calls
-// getKernelNameHelper(); a friend function that is only defined if
-// type KN was registered as a kernel name type in some translation unit
-// using the KernelRegistrar class template above. As with KernelRegistrar,
-// the call to getKernelNameHelper() requires ADL, so an instance of this
-// class is passed as an argument which is convertible to the kernel
-// name wrapper type.
-template <typename KN> struct KernelRegistry {
-#if defined(__GNUC__) && !defined(__clang__)
-#pragma GCC diagnostic push
-#pragma GCC diagnostic ignored "-Wnon-template-friend"
-#endif
-#if defined(_WIN32)
-  __declspec(dllexport)
-#endif
-  friend const char *getKernelNameHelper(KernelIdentity<KN>) noexcept;
-#if defined(__GNUC__) && !defined(__clang__)
-#pragma GCC diagnostic pop
-#endif
-  operator KernelIdentity<KN>() const noexcept { return KernelIdentity<KN>{}; }
-  static const char *getKernelName() noexcept {
-    return getKernelNameHelper(KernelRegistry<KN>{});
-  }
-};
-
 } // namespace detail
 } // namespace _V1
 } // namespace sycl

sycl/include/sycl/handler.hpp

@@ -799,8 +799,7 @@ class __SYCL_EXPORT handler {
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     constexpr auto Info = detail::CompileTimeKernelInfo<NameT>;
-    detail::KernelRegistrar<NameT,
-                            detail::KernelInfo<NameT>>::registerKernelName();
+
 #ifndef __SYCL_DEVICE_ONLY__
     throwIfActionIsCreated();
     throwOnKernelParameterMisuse(Info);
@@ -962,8 +961,6 @@ class __SYCL_EXPORT handler {
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     (void)Props;
     constexpr auto Info = detail::CompileTimeKernelInfo<NameT>;
-    detail::KernelRegistrar<NameT,
-                            detail::KernelInfo<NameT>>::registerKernelName();
     detail::KernelWrapper<WrapAsVal, NameT, KernelType, ElementType,
                           PropertiesT>::wrap(KernelFunc);
 

sycl/include/sycl/kernel_bundle.hpp

@@ -602,8 +602,10 @@ __SYCL_EXPORT kernel_id get_kernel_id_impl(string_view KernelName);
 
 /// \returns the kernel_id associated with the KernelName
 template <typename KernelName> kernel_id get_kernel_id() {
+  // FIXME: This must fail at link-time if KernelName not in any available
+  // translation units.
   return detail::get_kernel_id_impl(
-      std::string_view(detail::KernelRegistry<KernelName>::getKernelName()));
+      detail::CompileTimeKernelInfo<KernelName>.Name);
 }
 
 /// \returns a vector with all kernel_id's defined in the application

sycl/include/sycl/queue.hpp

@@ -3916,8 +3916,7 @@ auto submit_kernel_direct(const queue &Queue,
   detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
 
   using KernelType = std::decay_t<KernelTypeUniversalRef>;
-  detail::KernelRegistrar<KernelName,
-                          detail::KernelInfo<KernelName>>::registerKernelName();
+
   detail::KernelWrapper<WrapAs, KernelName, KernelType, LambdaArgType,
                         PropertiesT>::wrap(KernelFunc);
 

sycl/source/detail/device_kernel_info.cpp

@@ -50,17 +50,7 @@ void DeviceKernelInfo::setCompileTimeInfoIfNeeded(
   if (!isCompileTimeInfoSet())
     CompileTimeKernelInfoTy::operator=(Info);
   assert(isCompileTimeInfoSet());
-// FIXME On Windows, if we have a duplicate kernel name in multiple DSOs, the
-// full equality assertion can be violated. In this case, we would still be
-// reusing the kernel associated with this device kernel info (on Linux or
-// Windows), and while we could issue a diagnostic for the Windows case here,
-// we don't in order to maintain parity with Linux behavior, which happens
-// to work if the kernels are identical.
-#ifdef _WIN32
-  assert(std::string_view{Info.Name} == std::string_view{this->Name});
-#else
   assert(Info == *this);
-#endif
 }
 
 void DeviceKernelInfo::setImplicitLocalArgPos(int Pos) {

sycl/unittests/Extensions/RootGroup.cpp

@@ -43,8 +43,6 @@ struct KernelInfo<QueryKernel> : public unittest::MockKernelInfoBase {
 } // namespace _V1
 } // namespace sycl
 
-template void sycl::unittest::registerKernelNames<QueryKernel>();
-
 static sycl::unittest::MockDeviceImage Img =
     sycl::unittest::generateDefaultImage({"QueryKernel"});
 static const sycl::unittest::MockDeviceImageArray<1> ImgArray{&Img};

sycl/unittests/SYCL2020/KernelID.cpp

@@ -45,8 +45,6 @@ static sycl::unittest::MockDeviceImage Imgs[2] = {
     sycl::unittest::generateDefaultImage({"KernelID_TestKernel2"})};
 static sycl::unittest::MockDeviceImageArray<2> ImgArray{Imgs};
 
-template void
-sycl::unittest::registerKernelNames<TestKernel1, TestKernel2, TestKernel3>();
 TEST(KernelID, AllProgramKernelIds) {
   std::vector<sycl::kernel_id> AllKernelIDs = sycl::get_kernel_ids();
 
@@ -209,3 +207,18 @@ TEST(KernelID, HasKernelTemplated) {
   EXPECT_FALSE(InputBundle1.has_kernel<TestKernel2>());
   EXPECT_TRUE(InputBundle1.has_kernel<TestKernel3>());
 }
+
+TEST(KernelID, GetKernelIDInvalidKernelName) {
+  sycl::unittest::UrMock<> Mock;
+  sycl::platform Plt = sycl::platform();
+
+  try {
+    sycl::get_kernel_id<class NotAKernel>();
+    FAIL() << "Expected an exception";
+  } catch (sycl::exception const &e) {
+    EXPECT_TRUE(e.code() == sycl::errc::runtime);
+    EXPECT_EQ(std::string("No kernel found with the specified name"), e.what());
+  } catch (...) {
+    FAIL() << "Expected sycl::exception";
+  }
+}