[SYCL] Cache kernel_id and implement sycl::get_kernel_ids (#4379)

These changes implement the sycl::get_kernel_ids free function and
introduces caching of kernel identifiers in the program manager. The
caching of kernel identifiers is intended to ensure that each
identifier is only created once and copies of the identifier is passed
around, as to ensure equality in accordance with the SYCL 2020
specification.

NOTE: This is a non-breaking ABI change.

Signed-off-by: Steffen Larsen <[email protected]>

Author: steffenlarsen

sycl/include/CL/sycl/kernel_bundle.hpp

@@ -57,8 +57,6 @@ class __SYCL_EXPORT kernel_id {
 
   template <class T>
   friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
-
-  template <typename KernelName> friend kernel_id get_kernel_id();
 };
 
 namespace detail {
@@ -343,14 +341,20 @@ class kernel_bundle : public detail::kernel_bundle_plain {
 // get_kernel_id API
 /////////////////////////
 
+namespace detail {
+// Internal non-template versions of get_kernel_id API which is used by public
+// onces
+__SYCL_EXPORT kernel_id get_kernel_id_impl(std::string KernelName);
+} // namespace detail
+
 /// \returns the kernel_id associated with the KernelName
 template <typename KernelName> kernel_id get_kernel_id() {
   using KI = sycl::detail::KernelInfo<KernelName>;
-  return sycl::kernel_id(KI::getName());
+  return detail::get_kernel_id_impl(KI::getName());
 }
 
 /// \returns a vector with all kernel_id's defined in the application
-std::vector<kernel_id> get_kernel_ids();
+__SYCL_EXPORT std::vector<kernel_id> get_kernel_ids();
 
 /////////////////////////
 // get_kernel_bundle API

sycl/source/detail/program_manager/program_manager.cpp

@@ -1037,11 +1037,21 @@ void ProgramManager::addImages(pi_device_binaries DeviceBinary) {
       }
       // ... or create the set first if it hasn't been
       KernelSetId KSId = getNextKernelSetId();
-      for (_pi_offload_entry EntriesIt = EntriesB; EntriesIt != EntriesE;
-           ++EntriesIt) {
-        auto Result = KSIdMap.insert(std::make_pair(EntriesIt->name, KSId));
-        (void)Result;
-        assert(Result.second && "Kernel sets are not disjoint");
+      {
+        std::lock_guard<std::mutex> KernelIDsGuard(m_KernelIDsMutex);
+        for (_pi_offload_entry EntriesIt = EntriesB; EntriesIt != EntriesE;
+             ++EntriesIt) {
+          auto Result = KSIdMap.insert(std::make_pair(EntriesIt->name, KSId));
+          (void)Result;
+          assert(Result.second && "Kernel sets are not disjoint");
+          // ... 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);
+          sycl::kernel_id KernelID =
+              detail::createSyclObjFromImpl<sycl::kernel_id>(KernelIDImpl);
+          m_KernelIDs.insert(
+              std::make_pair(EntriesIt->name, std::move(KernelID)));
+        }
       }
       m_DeviceImages[KSId].reset(new std::vector<RTDeviceBinaryImageUPtr>());
       m_DeviceImages[KSId]->push_back(std::move(Img));
@@ -1266,6 +1276,25 @@ static bool compatibleWithDevice(RTDeviceBinaryImage *BinImage,
   return (0 == SuitableImageID);
 }
 
+kernel_id ProgramManager::getSYCLKernelID(const std::string &KernelName) {
+  std::lock_guard<std::mutex> KernelIDsGuard(m_KernelIDsMutex);
+
+  auto KernelID = m_KernelIDs.find(KernelName);
+  assert(KernelID != m_KernelIDs.end() && "Kernel ID missing");
+  return KernelID->second;
+}
+
+std::vector<kernel_id> ProgramManager::getAllSYCLKernelIDs() {
+  std::lock_guard<std::mutex> KernelIDsGuard(m_KernelIDsMutex);
+
+  std::vector<sycl::kernel_id> AllKernelIDs;
+  AllKernelIDs.reserve(m_KernelIDs.size());
+  for (std::pair<std::string, kernel_id> KernelID : m_KernelIDs) {
+    AllKernelIDs.push_back(KernelID.second);
+  }
+  return AllKernelIDs;
+}
+
 std::vector<device_image_plain>
 ProgramManager::getSYCLDeviceImagesWithCompatibleState(
     const context &Ctx, const std::vector<device> &Devs,
@@ -1317,14 +1346,15 @@ ProgramManager::getSYCLDeviceImagesWithCompatibleState(
       // Collect kernel names for the image
       pi_device_binary DevBin =
           const_cast<pi_device_binary>(&BinImage->getRawData());
-      for (_pi_offload_entry EntriesIt = DevBin->EntriesBegin;
-           EntriesIt != DevBin->EntriesEnd; ++EntriesIt) {
-
-        std::shared_ptr<detail::kernel_id_impl> KernelIDImpl =
-            std::make_shared<detail::kernel_id_impl>(EntriesIt->name);
-
-        KernelIDs.push_back(
-            detail::createSyclObjFromImpl<sycl::kernel_id>(KernelIDImpl));
+      {
+        std::lock_guard<std::mutex> KernelIDsGuard(m_KernelIDsMutex);
+        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");
+          KernelIDs.push_back(KernelID->second);
+        }
       }
       // 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

@@ -153,6 +153,14 @@ class ProgramManager {
                                            const std::string &KernelName,
                                            bool KnownProgram);
 
+  // The function returns the unique SYCL kernel identifier associated with a
+  // kernel name.
+  kernel_id getSYCLKernelID(const std::string &KernelName);
+
+  // The function returns a vector containing all unique SYCL kernel identifiers
+  // in SYCL device images.
+  std::vector<kernel_id> getAllSYCLKernelIDs();
+
   // The function returns a vector of SYCL device images that are compiled with
   // the required state and at least one device from the passed list of devices.
   std::vector<device_image_plain>
@@ -272,6 +280,18 @@ class ProgramManager {
   /// Access must be guarded by the \ref Sync::getGlobalLock()
   std::unordered_map<OSModuleHandle, KernelSetId> m_OSModuleKernelSets;
 
+  /// 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
+  std::unordered_map<std::string, kernel_id> m_KernelIDs;
+
+  /// Protects kernel ID cache.
+  /// NOTE: This may be acquired while \ref Sync::getGlobalLock() is held so to
+  /// avoid deadlocks care must be taken not to acquire
+  /// \ref Sync::getGlobalLock() while holding this mutex.
+  std::mutex m_KernelIDsMutex;
+
   // 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/source/kernel_bundle.cpp

@@ -110,9 +110,13 @@ bool kernel_bundle_plain::is_specialization_constant_set(
   return impl->is_specialization_constant_set(SpecName);
 }
 
-////////////////////////////
-///// free functions
-///////////////////////////
+//////////////////////////////////
+///// sycl::detail free functions
+//////////////////////////////////
+
+kernel_id get_kernel_id_impl(std::string KernelName) {
+  return detail::ProgramManager::getInstance().getSYCLKernelID(KernelName);
+}
 
 detail::KernelBundleImplPtr
 get_kernel_bundle_impl(const context &Ctx, const std::vector<device> &Devs,
@@ -263,5 +267,13 @@ std::vector<sycl::device> find_device_intersection(
 
 } // namespace detail
 
+//////////////////////////
+///// sycl free functions
+//////////////////////////
+
+std::vector<kernel_id> get_kernel_ids() {
+  return detail::ProgramManager::getInstance().getAllSYCLKernelIDs();
+}
+
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/test/abi/sycl_symbols_linux.dump

@@ -3618,6 +3618,7 @@ _ZN2cl4sycl13sycl_categoryEv
 _ZN2cl4sycl14exception_list5ClearEv
 _ZN2cl4sycl14exception_list8PushBackEONSt15__exception_ptr13exception_ptrE
 _ZN2cl4sycl14exception_list8PushBackERKNSt15__exception_ptr13exception_ptrE
+_ZN2cl4sycl14get_kernel_idsEv
 _ZN2cl4sycl15make_error_codeENS0_4errcE
 _ZN2cl4sycl16get_pointer_typeEPKvRKNS0_7contextE
 _ZN2cl4sycl18aligned_alloc_hostEmmRKNS0_5queueE
@@ -3818,6 +3819,7 @@ _ZN2cl4sycl6detail17HostProfilingInfo3endEv
 _ZN2cl4sycl6detail17HostProfilingInfo5startEv
 _ZN2cl4sycl6detail18convertChannelTypeE22_pi_image_channel_type
 _ZN2cl4sycl6detail18convertChannelTypeENS0_18image_channel_typeE
+_ZN2cl4sycl6detail18get_kernel_id_implENSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
 _ZN2cl4sycl6detail18make_kernel_bundleEmRKNS0_7contextENS0_12bundle_stateENS0_7backendE
 _ZN2cl4sycl6detail18stringifyErrorCodeEi
 _ZN2cl4sycl6detail19convertChannelOrderE23_pi_image_channel_order

sycl/test/abi/sycl_symbols_windows.dump

@@ -2133,7 +2133,9 @@
 ?get_kernel_bundle_impl@detail@sycl@cl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@sycl@cl@@@std@@AEBVcontext@23@AEBV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@5@AEBV?$vector@Vkernel_id@sycl@cl@@V?$allocator@Vkernel_id@sycl@cl@@@std@@@5@W4bundle_state@23@@Z
 ?get_kernel_bundle_impl@detail@sycl@cl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@sycl@cl@@@std@@AEBVcontext@23@AEBV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@5@W4bundle_state@23@@Z
 ?get_kernel_bundle_impl@detail@sycl@cl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@sycl@cl@@@std@@AEBVcontext@23@AEBV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@5@W4bundle_state@23@AEBV?$function@$$A6A_NAEBV?$shared_ptr@Vdevice_image_impl@detail@sycl@cl@@@std@@@Z@5@@Z
+?get_kernel_id_impl@detail@sycl@cl@@YA?AVkernel_id@23@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@Z
 ?get_kernel_ids@kernel_bundle_plain@detail@sycl@cl@@QEBA?AV?$vector@Vkernel_id@sycl@cl@@V?$allocator@Vkernel_id@sycl@cl@@@std@@@std@@XZ
+?get_kernel_ids@sycl@cl@@YA?AV?$vector@Vkernel_id@sycl@cl@@V?$allocator@Vkernel_id@sycl@cl@@@std@@@std@@XZ
 ?get_link_options@program@sycl@cl@@QEBA?AV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@XZ
 ?get_max_statement_size@stream@sycl@cl@@QEBA_KXZ
 ?get_max_statement_size@stream_impl@detail@sycl@cl@@QEBA_KXZ

sycl/unittests/SYCL2020/CMakeLists.txt

@@ -6,5 +6,6 @@ add_sycl_unittest(SYCL2020Tests OBJECT
   GetNativeOpenCL.cpp
   SpecConstDefaultValues.cpp
   KernelBundle.cpp
+  KernelID.cpp
 )