[SYCL] Refactor kernel name based cache approach

Author: sergey-semenov

sycl/include/sycl/detail/get_kernel_name_based_data.hpp

@@ -0,0 +1,38 @@
+//==--------------------- get_kernel_name_based_data.hpp -------------------==//
+//
+// 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
+
+#include <sycl/detail/export.hpp>
+#include <sycl/detail/kernel_name_str_t.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+namespace detail {
+
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+class KernelNameBasedCacheT;
+__SYCL_EXPORT KernelNameBasedCacheT *createKernelNameBasedCache();
+#endif
+
+class KernelNameBasedData;
+
+__SYCL_EXPORT KernelNameBasedData *
+getKernelNameBasedDataImpl(detail::ABINeutralKernelNameStrRefT KernelName);
+
+// Retrieves and caches a data pointer to avoid kernel name based lookup
+// overhead.
+template <typename KernelNameT>
+KernelNameBasedData *
+getKernelNameBasedData(detail::ABINeutralKernelNameStrRefT KernelName) {
+  static KernelNameBasedData *Instance = getKernelNameBasedDataImpl(KernelName);
+  return Instance;
+}
+
+} // namespace detail
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/detail/kernel_name_based_cache.hpp

@@ -18,10 +18,12 @@ namespace detail {
 using KernelNameStrT = std::string_view;
 using KernelNameStrRefT = std::string_view;
 using ABINeutralKernelNameStrT = detail::string_view;
+using ABINeutralKernelNameStrRefT = detail::string_view;
 #else
 using KernelNameStrT = std::string;
 using KernelNameStrRefT = const std::string &;
 using ABINeutralKernelNameStrT = detail::string;
+using ABINeutralKernelNameStrRefT = const detail::string &;
 #endif
 
 inline KernelNameStrT toKernelNameStrT(const ABINeutralKernelNameStrT &str) {

sycl/include/sycl/detail/kernel_name_str_t.hpp

@@ -14,11 +14,11 @@
 #include <sycl/detail/common.hpp>
 #include <sycl/detail/defines_elementary.hpp>
 #include <sycl/detail/export.hpp>
+#include <sycl/detail/get_kernel_name_based_data.hpp>
 #include <sycl/detail/id_queries_fit_in_int.hpp>
 #include <sycl/detail/impl_utils.hpp>
 #include <sycl/detail/kernel_desc.hpp>
 #include <sycl/detail/kernel_launch_helper.hpp>
-#include <sycl/detail/kernel_name_based_cache.hpp>
 #include <sycl/detail/kernel_name_str_t.hpp>
 #include <sycl/detail/reduction_forward.hpp>
 #include <sycl/detail/string.hpp>
@@ -863,14 +863,15 @@ class __SYCL_EXPORT handler {
       constexpr std::string_view KernelNameStr =
           detail::getKernelName<KernelName>();
       MKernelName = KernelNameStr;
+      setKernelNameBasedDataPtr(
+          detail::getKernelNameBasedData<KernelName>(KernelNameStr));
     } else {
       // In case w/o the integration header it is necessary to process
       // accessors from the list(which are associated with this handler) as
       // arguments. We must copy the associated accessors as they are checked
       // later during finalize.
       setArgsToAssociatedAccessors();
     }
-    setKernelNameBasedCachePtr(detail::getKernelNameBasedCache<KernelName>());
 
     // If the kernel lambda is callable with a kernel_handler argument, manifest
     // the associated kernel handler.
@@ -3685,8 +3686,12 @@ class __SYCL_EXPORT handler {
       sycl::handler &h, size_t size,
       const ext::oneapi::experimental::memory_pool &pool);
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   void setKernelNameBasedCachePtr(
       detail::KernelNameBasedCacheT *KernelNameBasedCachePtr);
+#endif
+  void setKernelNameBasedDataPtr(
+      detail::KernelNameBasedData *KernelNameBasedDataPtr);
 
   queue getQueue();
 

sycl/include/sycl/handler.hpp

@@ -288,7 +288,8 @@ set(SYCL_COMMON_SOURCES
     "detail/kernel_compiler/kernel_compiler_opencl.cpp"
     "detail/kernel_compiler/kernel_compiler_sycl.cpp"
     "detail/kernel_impl.cpp"
-    "detail/kernel_name_based_cache.cpp"
+    "detail/get_kernel_name_based_data.cpp"
+    "detail/kernel_name_based_data.cpp"
     "detail/kernel_program_cache.cpp"
     "detail/memory_export.cpp"
     "detail/memory_manager.cpp"

sycl/source/CMakeLists.txt

@@ -254,7 +254,7 @@ class CGExecKernel : public CG {
   std::shared_ptr<detail::kernel_bundle_impl> MKernelBundle;
   std::vector<ArgDesc> MArgs;
   KernelNameStrT MKernelName;
-  KernelNameBasedCacheT *MKernelNameBasedCachePtr;
+  KernelNameBasedData *MKernelNameBasedDataPtr;
   std::vector<std::shared_ptr<detail::stream_impl>> MStreams;
   std::vector<std::shared_ptr<const void>> MAuxiliaryResources;
   /// Used to implement ext_oneapi_graph dynamic_command_group. Stores the list
@@ -270,7 +270,7 @@ class CGExecKernel : public CG {
                std::shared_ptr<detail::kernel_bundle_impl> KernelBundle,
                CG::StorageInitHelper CGData, std::vector<ArgDesc> Args,
                KernelNameStrT KernelName,
-               KernelNameBasedCacheT *KernelNameBasedCachePtr,
+               KernelNameBasedData *KernelNameBasedDataPtr,
                std::vector<std::shared_ptr<detail::stream_impl>> Streams,
                std::vector<std::shared_ptr<const void>> AuxiliaryResources,
                CGType Type, ur_kernel_cache_config_t KernelCacheConfig,
@@ -280,7 +280,7 @@ class CGExecKernel : public CG {
         MHostKernel(std::move(HKernel)), MSyclKernel(std::move(SyclKernel)),
         MKernelBundle(std::move(KernelBundle)), MArgs(std::move(Args)),
         MKernelName(std::move(KernelName)),
-        MKernelNameBasedCachePtr(KernelNameBasedCachePtr),
+        MKernelNameBasedDataPtr(KernelNameBasedDataPtr),
         MStreams(std::move(Streams)),
         MAuxiliaryResources(std::move(AuxiliaryResources)),
         MAlternativeKernels{}, MKernelCacheConfig(std::move(KernelCacheConfig)),

sycl/source/detail/cg.hpp

@@ -1,4 +1,4 @@
-//==--------------------- kernel_name_based_cache.cpp ----------------------==//
+//==-------------------- get_kernel_name_based_data.cpp --------------------==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -7,15 +7,24 @@
 //===----------------------------------------------------------------------===//
 
 #include <detail/global_handler.hpp>
-#include <sycl/detail/kernel_name_based_cache.hpp>
+#include <detail/program_manager/program_manager.hpp>
+#include <sycl/detail/get_kernel_name_based_data.hpp>
 
 namespace sycl {
 inline namespace _V1 {
 namespace detail {
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 KernelNameBasedCacheT *createKernelNameBasedCache() {
   return GlobalHandler::instance().createKernelNameBasedCache();
 }
+#endif
+
+KernelNameBasedData *
+getKernelNameBasedDataImpl(detail::ABINeutralKernelNameStrRefT KernelName) {
+  return ProgramManager::getInstance().getOrCreateKernelNameBasedData(
+      KernelName.data());
+}
 
 } // namespace detail
 } // namespace _V1

sycl/source/detail/kernel_name_based_cache.cpp renamed to sycl/source/detail/get_kernel_name_based_data.cpp

@@ -14,7 +14,7 @@
 #include <detail/adapter_impl.hpp>
 #include <detail/config.hpp>
 #include <detail/global_handler.hpp>
-#include <detail/kernel_name_based_cache_t.hpp>
+#include <detail/kernel_name_based_data.hpp>
 #include <detail/platform_impl.hpp>
 #include <detail/program_manager/program_manager.hpp>
 #include <detail/scheduler/scheduler.hpp>
@@ -249,12 +249,15 @@ ThreadPool &GlobalHandler::getHostTaskThreadPool() {
   return TP;
 }
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 KernelNameBasedCacheT *GlobalHandler::createKernelNameBasedCache() {
-  static std::deque<KernelNameBasedCacheT> &KernelNameBasedCaches =
-      getOrCreate(MKernelNameBasedCaches);
-  LockGuard LG{MKernelNameBasedCaches.Lock};
-  return &KernelNameBasedCaches.emplace_back();
+  static std::deque<KernelNameBasedData> &KernelNameBasedDataStorage =
+      getOrCreate(MKernelNameBasedDataStorage);
+  LockGuard LG{MKernelNameBasedDataStorage.Lock};
+  return reinterpret_cast<KernelNameBasedCacheT *>(
+      &KernelNameBasedDataStorage.emplace_back());
 }
+#endif
 
 void GlobalHandler::releaseDefaultContexts() {
   // Release shared-pointers to SYCL objects.
@@ -390,9 +393,11 @@ void shutdown_late() {
   Handler->MScheduler.Inst.reset(nullptr);
   Handler->MProgramManager.Inst.reset(nullptr);
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   // Cache stores handles to the adapter, so clear it before
   // releasing adapters.
-  Handler->MKernelNameBasedCaches.Inst.reset(nullptr);
+  Handler->MKernelNameBasedDataStorage.Inst.reset(nullptr);
+#endif
 
   // Clear the adapters and reset the instance if it was there.
   Handler->unloadAdapters();

sycl/source/detail/global_handler.cpp

@@ -11,7 +11,9 @@
 #include <sycl/detail/spinlock.hpp>
 #include <sycl/detail/util.hpp>
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 #include <deque>
+#endif
 #include <memory>
 #include <unordered_map>
 
@@ -27,7 +29,10 @@ class adapter_impl;
 class ods_target_list;
 class XPTIRegistry;
 class ThreadPool;
-struct KernelNameBasedCacheT;
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+class KernelNameBasedCacheT;
+class KernelNameBasedData;
+#endif
 
 /// Wrapper class for global data structures with non-trivial destructors.
 ///
@@ -73,7 +78,9 @@ class GlobalHandler {
   ods_target_list &getOneapiDeviceSelectorTargets(const std::string &InitValue);
   XPTIRegistry &getXPTIRegistry();
   ThreadPool &getHostTaskThreadPool();
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   KernelNameBasedCacheT *createKernelNameBasedCache();
+#endif
   static void registerStaticVarShutdownHandler();
 
   bool isOkToDefer() const;
@@ -130,7 +137,9 @@ class GlobalHandler {
   InstWithLock<XPTIRegistry> MXPTIRegistry;
   // Thread pool for host task and event callbacks execution
   InstWithLock<ThreadPool> MHostTaskThreadPool;
-  InstWithLock<std::deque<KernelNameBasedCacheT>> MKernelNameBasedCaches;
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+  InstWithLock<std::deque<KernelNameBasedData>> MKernelNameBasedDataStorage;
+#endif
 };
 } // namespace detail
 } // namespace _V1

sycl/source/detail/global_handler.hpp

@@ -741,7 +741,7 @@ ur_exp_command_buffer_sync_point_t exec_graph_impl::enqueueNodeDirect(
                                         CGExec->MLine, CGExec->MColumn);
     std::tie(CmdTraceEvent, InstanceID) = emitKernelInstrumentationData(
         StreamID, CGExec->MSyclKernel, CodeLoc, CGExec->MIsTopCodeLoc,
-        CGExec->MKernelName.data(), CGExec->MKernelNameBasedCachePtr, nullptr,
+        CGExec->MKernelName.data(), CGExec->MKernelNameBasedDataPtr, nullptr,
         CGExec->MNDRDesc, CGExec->MKernelBundle.get(), CGExec->MArgs);
     if (CmdTraceEvent)
       sycl::detail::emitInstrumentationGeneral(
@@ -1573,9 +1573,10 @@ void exec_graph_impl::populateURKernelUpdateStructs(
     UrKernel = SyclKernelImpl->getHandleRef();
     EliminatedArgMask = SyclKernelImpl->getKernelArgMask();
   } else {
+    assert(ExecCG.MKernelNameBasedDataPtr);
     BundleObjs = sycl::detail::ProgramManager::getInstance().getOrCreateKernel(
         ContextImpl, DeviceImpl, ExecCG.MKernelName,
-        ExecCG.MKernelNameBasedCachePtr);
+        *ExecCG.MKernelNameBasedDataPtr);
     UrKernel = BundleObjs->MKernelHandle;
     EliminatedArgMask = BundleObjs->MKernelArgMask;
   }