[SYCL] Hide SYCL service kernels (#4519)

The SYCL runtime may in places benefit from being able to define
"service kernels" to use behind-the-scenes work on devices. These
kernels may use the same interface as other SYCL kernels, but as such
will be handled no different than user-defined kernels.

These changes makes a distinction between service kernels and
user-defined kernels by placing service kernels in a common namespace
(`cl::sycl::detail::__sycl_service_kernel__`). The program manager will
not grant unique kernel IDs to service kernels. This prevents service
kernels from being visible to the user through kernel_bundle related
interfaces.

Consequently the runtime may create device images that do not contain
any kernel IDs. As such, `sycl::has_kernel_bundle` will now discount
any device images that have no kernel IDs.

Signed-off-by: Steffen Larsen <steffen.larsen@intel.com>

Author: steffenlarsen

sycl/include/CL/sycl/detail/service_kernel_names.hpp

@@ -0,0 +1,21 @@
+//==-------- service_kernels.hpp - SYCL service kernel name types ----------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {
+namespace __sycl_service_kernel__ {
+
+class AssertInfoCopier;
+
+} // namespace __sycl_service_kernel__
+} // namespace detail
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/queue.hpp

@@ -12,6 +12,7 @@
 #include <CL/sycl/detail/assert_happened.hpp>
 #include <CL/sycl/detail/common.hpp>
 #include <CL/sycl/detail/export.hpp>
+#include <CL/sycl/detail/service_kernel_names.hpp>
 #include <CL/sycl/device.hpp>
 #include <CL/sycl/device_selector.hpp>
 #include <CL/sycl/event.hpp>
@@ -79,7 +80,6 @@ class queue;
 namespace detail {
 class queue_impl;
 #if __SYCL_USE_FALLBACK_ASSERT
-class AssertInfoCopier;
 static event submitAssertCapture(queue &, event &, queue *,
                                  const detail::code_location &);
 #endif
@@ -1172,7 +1172,7 @@ event submitAssertCapture(queue &Self, event &Event, queue *SecondaryQueue,
 
     auto Acc = Buffer.get_access<access::mode::write>(CGH);
 
-    CGH.single_task<AssertInfoCopier>([Acc] {
+    CGH.single_task<__sycl_service_kernel__::AssertInfoCopier>([Acc] {
 #ifdef __SYCL_DEVICE_ONLY__
       __devicelib_assert_read(&Acc[0]);
 #else

sycl/source/detail/program_manager/program_manager.cpp

@@ -1064,6 +1064,16 @@ void ProgramManager::addImages(pi_device_binaries DeviceBinary) {
           auto Result = KSIdMap.insert(std::make_pair(EntriesIt->name, KSId));
           (void)Result;
           assert(Result.second && "Kernel sets are not disjoint");
+
+          // Skip creating unique kernel ID if it is a service kernel.
+          // SYCL service kernels are identified by having
+          // __sycl_service_kernel__ in the mangled name, primarily as part of
+          // the namespace of the name type.
+          if (std::strstr(EntriesIt->name, "__sycl_service_kernel__")) {
+            m_ServiceKernels.insert(EntriesIt->name);
+            continue;
+          }
+
           // ... and create a unique kernel ID for the entry
           std::shared_ptr<detail::kernel_id_impl> KernelIDImpl =
               std::make_shared<detail::kernel_id_impl>(EntriesIt->name);
@@ -1352,7 +1362,6 @@ ProgramManager::getSYCLDeviceImagesWithCompatibleState(
       if (!compatibleWithDevice(BinImage, Dev))
         continue;
 
-      // TODO: Cache kernel_ids
       std::vector<sycl::kernel_id> KernelIDs;
       // Collect kernel names for the image
       pi_device_binary DevBin =
@@ -1362,11 +1371,23 @@ ProgramManager::getSYCLDeviceImagesWithCompatibleState(
         for (_pi_offload_entry EntriesIt = DevBin->EntriesBegin;
              EntriesIt != DevBin->EntriesEnd; ++EntriesIt) {
           auto KernelID = m_KernelIDs.find(EntriesIt->name);
-          assert(KernelID != m_KernelIDs.end() &&
-                 "Kernel ID in device binary missing from cache");
+
+          if (KernelID == m_KernelIDs.end()) {
+            // Service kernels do not have kernel IDs
+            assert(m_ServiceKernels.find(EntriesIt->name) !=
+                       m_ServiceKernels.end() &&
+                   "Kernel ID in device binary missing from cache");
+            continue;
+          }
+
           KernelIDs.push_back(KernelID->second);
         }
       }
+
+      // If the image does not contain any non-service kernels we can skip it.
+      if (KernelIDs.empty())
+        continue;
+
       // device_image_impl expects kernel ids to be sorted for fast search
       std::sort(KernelIDs.begin(), KernelIDs.end(), LessByNameComp{});
 

sycl/source/detail/program_manager/program_manager.hpp

@@ -23,6 +23,7 @@
 #include <memory>
 #include <set>
 #include <unordered_map>
+#include <unordered_set>
 #include <vector>
 
 // +++ Entry points referenced by the offload wrapper object {
@@ -281,7 +282,7 @@ class ProgramManager {
   /// Maps names of kernels to their unique kernel IDs.
   /// TODO: Use std::unordered_set with transparent hash and equality functions
   ///       when C++20 is enabled for the runtime library.
-  /// Access must be guarded by the m_KernelIDsMutex mutex
+  /// Access must be guarded by the m_KernelIDsMutex mutex.
   std::unordered_map<std::string, kernel_id> m_KernelIDs;
 
   /// Protects kernel ID cache.
@@ -290,6 +291,14 @@ class ProgramManager {
   /// \ref Sync::getGlobalLock() while holding this mutex.
   std::mutex m_KernelIDsMutex;
 
+  /// Caches all found service kernels to expedite future checks. A SYCL service
+  /// kernel is a kernel that has not been defined by the user but is instead
+  /// generated by the SYCL runtime. Service kernel name types must be declared
+  /// in the sycl::detail::__sycl_service_kernel__ namespace which is
+  /// exclusively used for this purpose.
+  /// Access must be guarded by the m_KernelIDsMutex mutex.
+  std::unordered_set<std::string> m_ServiceKernels;
+
   // Keeps track of pi_program to image correspondence. Needed for:
   // - knowing which specialization constants are used in the program and
   //   injecting their current values before compiling the SPIR-V; the binary

sycl/unittests/SYCL2020/KernelID.cpp

@@ -17,6 +17,7 @@
 class TestKernel1;
 class TestKernel2;
 class TestKernel3;
+class ServiceKernel1;
 
 __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
@@ -57,6 +58,19 @@ template <> struct KernelInfo<TestKernel3> {
   static constexpr bool callsAnyThisFreeFunction() { return false; }
 };
 
+template <> struct KernelInfo<ServiceKernel1> {
+  static constexpr unsigned getNumParams() { return 0; }
+  static const kernel_param_desc_t &getParamDesc(int) {
+    static kernel_param_desc_t Dummy;
+    return Dummy;
+  }
+  static constexpr const char *getName() {
+    return "_ZTSN2cl4sycl6detail23__sycl_service_kernel__14ServiceKernel1";
+  }
+  static constexpr bool isESIMD() { return false; }
+  static constexpr bool callsThisItem() { return false; }
+  static constexpr bool callsAnyThisFreeFunction() { return false; }
+};
 } // namespace detail
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)
@@ -84,7 +98,9 @@ generateDefaultImage(std::initializer_list<std::string> Kernels) {
 
 static sycl::unittest::PiImage Imgs[2] = {
     generateDefaultImage({"KernelID_TestKernel1", "KernelID_TestKernel3"}),
-    generateDefaultImage({"KernelID_TestKernel2"})};
+    generateDefaultImage(
+        {"KernelID_TestKernel2",
+         "_ZTSN2cl4sycl6detail23__sycl_service_kernel__14ServiceKernel1"})};
 static sycl::unittest::PiImageArray<2> ImgArray{Imgs};
 
 TEST(KernelID, AllProgramKernelIds) {
@@ -106,6 +122,20 @@ TEST(KernelID, AllProgramKernelIds) {
   }
 }
 
+TEST(KernelID, NoServiceKernelIds) {
+  const char *ServiceKernel1Name =
+      sycl::detail::KernelInfo<ServiceKernel1>::getName();
+
+  std::vector<sycl::kernel_id> AllKernelIDs = sycl::get_kernel_ids();
+
+  auto NoFoundServiceKernelID = std::none_of(
+      AllKernelIDs.begin(), AllKernelIDs.end(), [=](sycl::kernel_id KernelID) {
+        return strcmp(KernelID.get_name(), ServiceKernel1Name) == 0;
+      });
+
+  EXPECT_TRUE(NoFoundServiceKernelID);
+}
+
 TEST(KernelID, FreeKernelIDEqualsKernelBundleId) {
   sycl::platform Plt{sycl::default_selector()};
   if (Plt.is_host()) {

sycl/unittests/assert/assert.cpp

@@ -51,9 +51,10 @@ template <> struct KernelInfo<TestKernel> {
 static constexpr const kernel_param_desc_t Signatures[] = {
     {kernel_param_kind_t::kind_accessor, 4062, 0}};
 
-template <> struct KernelInfo<::sycl::detail::AssertInfoCopier> {
+template <>
+struct KernelInfo<::sycl::detail::__sycl_service_kernel__::AssertInfoCopier> {
   static constexpr const char *getName() {
-    return "_ZTSN2cl4sycl6detail16AssertInfoCopierE";
+    return "_ZTSN2cl4sycl6detail23__sycl_service_kernel__16AssertInfoCopierE";
   }
   static constexpr unsigned getNumParams() { return 1; }
   static constexpr const kernel_param_desc_t &getParamDesc(unsigned Idx) {
@@ -73,7 +74,7 @@ static sycl::unittest::PiImage generateDefaultImage() {
 
   static const std::string KernelName = "TestKernel";
   static const std::string CopierKernelName =
-      "_ZTSN2cl4sycl6detail16AssertInfoCopierE";
+      "_ZTSN2cl4sycl6detail23__sycl_service_kernel__16AssertInfoCopierE";
 
   PiPropertySet PropSet;
 
@@ -98,7 +99,7 @@ static sycl::unittest::PiImage generateCopierKernelImage() {
   using namespace sycl::unittest;
 
   static const std::string CopierKernelName =
-      "_ZTSN2cl4sycl6detail16AssertInfoCopierE";
+      "_ZTSN2cl4sycl6detail23__sycl_service_kernel__16AssertInfoCopierE";
 
   PiPropertySet PropSet;
 
@@ -396,3 +397,17 @@ TEST(Assert, TestPositive) {
   }
 #endif // _WIN32
 }
+
+TEST(Assert, TestAssertServiceKernelHidden) {
+  const char *AssertServiceKernelName = sycl::detail::KernelInfo<
+      sycl::detail::__sycl_service_kernel__::AssertInfoCopier>::getName();
+
+  std::vector<sycl::kernel_id> AllKernelIDs = sycl::get_kernel_ids();
+
+  auto NoFoundServiceKernelID = std::none_of(
+      AllKernelIDs.begin(), AllKernelIDs.end(), [=](sycl::kernel_id KernelID) {
+        return strcmp(KernelID.get_name(), AssertServiceKernelName) == 0;
+      });
+
+  EXPECT_TRUE(NoFoundServiceKernelID);
+}