[NFCI][SYCL] Switch to using Managed<ur_kernel_handle_t> (#19570)

aelovikov-intel · web-flow · commit 0dc91aa007d1 · 2025-07-25T08:01:05.000-07:00
Similar to what's been done with `ur_program_handle_t` in #19536 #19557
diff --git a/sycl/source/backend.cpp b/sycl/source/backend.cpp
@@ -321,7 +321,7 @@ kernel make_kernel(const context &TargetContext,
                    const kernel_bundle<bundle_state::executable> &KernelBundle,
                    ur_native_handle_t NativeHandle, bool KeepOwnership,
                    backend Backend) {
-  const auto &Adapter = getAdapter(Backend);
+  adapter_impl &Adapter = getAdapter(Backend);
   context_impl &ContextImpl = *getSyclObjImpl(TargetContext);
   kernel_bundle_impl &KernelBundleImpl = *getSyclObjImpl(KernelBundle);
 
@@ -347,7 +347,7 @@ kernel make_kernel(const context &TargetContext,
   }
 
   // Create UR kernel first.
-  ur_kernel_handle_t UrKernel = nullptr;
+  Managed<ur_kernel_handle_t> UrKernel{Adapter};
   ur_kernel_native_properties_t Properties{};
   Properties.stype = UR_STRUCTURE_TYPE_KERNEL_NATIVE_PROPERTIES;
   Properties.isNativeHandleOwned = !KeepOwnership;
@@ -359,8 +359,8 @@ kernel make_kernel(const context &TargetContext,
     __SYCL_OCL_CALL(clRetainKernel, ur::cast<cl_kernel>(NativeHandle));
 
   // Construct the SYCL queue from UR queue.
-  return detail::createSyclObjFromImpl<kernel>(
-      std::make_shared<kernel_impl>(UrKernel, ContextImpl, &KernelBundleImpl));
+  return detail::createSyclObjFromImpl<kernel>(std::make_shared<kernel_impl>(
+      std::move(UrKernel), ContextImpl, &KernelBundleImpl));
 }
 
 kernel make_kernel(ur_native_handle_t NativeHandle,
diff --git a/sycl/source/detail/adapter_impl.hpp b/sycl/source/detail/adapter_impl.hpp
@@ -243,10 +243,14 @@ template <typename URResource> class Managed {
   static constexpr auto Release = []() constexpr {
     if constexpr (std::is_same_v<URResource, ur_program_handle_t>)
       return UrApiKind::urProgramRelease;
+    if constexpr (std::is_same_v<URResource, ur_kernel_handle_t>)
+      return UrApiKind::urKernelRelease;
   }();
   static constexpr auto Retain = []() constexpr {
     if constexpr (std::is_same_v<URResource, ur_program_handle_t>)
       return UrApiKind::urProgramRetain;
+    if constexpr (std::is_same_v<URResource, ur_kernel_handle_t>)
+      return UrApiKind::urKernelRetain;
   }();
 
 public:
diff --git a/sycl/source/detail/device_image_impl.cpp b/sycl/source/detail/device_image_impl.cpp
@@ -34,27 +34,26 @@ std::shared_ptr<kernel_impl> device_image_impl::tryGetExtensionKernel(
       auto [UrKernel, CacheMutex, ArgMask] =
           PM.getOrCreateKernel(Context, AdjustedName,
                                /*PropList=*/{}, UrProgram);
-      return std::make_shared<kernel_impl>(UrKernel, *getSyclObjImpl(Context),
-                                           shared_from_this(), OwnerBundle,
-                                           ArgMask, UrProgram, CacheMutex);
+      return std::make_shared<kernel_impl>(
+          std::move(UrKernel), *getSyclObjImpl(Context), shared_from_this(),
+          OwnerBundle, ArgMask, UrProgram, CacheMutex);
     }
     return nullptr;
   }
 
   ur_program_handle_t UrProgram = get_ur_program();
   detail::adapter_impl &Adapter = getSyclObjImpl(Context)->getAdapter();
-  ur_kernel_handle_t UrKernel = nullptr;
+  Managed<ur_kernel_handle_t> UrKernel{Adapter};
   Adapter.call<UrApiKind::urKernelCreate>(UrProgram, AdjustedName.c_str(),
                                           &UrKernel);
-  // Kernel created by urKernelCreate is implicitly retained.
 
   const KernelArgMask *ArgMask = nullptr;
   if (auto ArgMaskIt = MEliminatedKernelArgMasks.find(AdjustedName);
       ArgMaskIt != MEliminatedKernelArgMasks.end())
     ArgMask = &ArgMaskIt->second;
 
   return std::make_shared<kernel_impl>(
-      UrKernel, *detail::getSyclObjImpl(Context), shared_from_this(),
+      std::move(UrKernel), *detail::getSyclObjImpl(Context), shared_from_this(),
       OwnerBundle, ArgMask, UrProgram, /*CacheMutex=*/nullptr);
 }
 
diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp
@@ -983,8 +983,9 @@ class kernel_bundle_impl
             SelectedImage->get_ur_program());
 
     return std::make_shared<kernel_impl>(
-        Kernel, *detail::getSyclObjImpl(MContext), std::move(SelectedImage),
-        *this, ArgMask, SelectedImage->get_ur_program(), CacheMutex);
+        std::move(Kernel), *detail::getSyclObjImpl(MContext),
+        std::move(SelectedImage), *this, ArgMask,
+        SelectedImage->get_ur_program(), CacheMutex);
   }
 
   std::shared_ptr<kernel_impl>
diff --git a/sycl/source/detail/kernel_impl.cpp b/sycl/source/detail/kernel_impl.cpp
@@ -16,11 +16,12 @@ namespace sycl {
 inline namespace _V1 {
 namespace detail {
 
-kernel_impl::kernel_impl(ur_kernel_handle_t Kernel, context_impl &Context,
+kernel_impl::kernel_impl(Managed<ur_kernel_handle_t> &&Kernel,
+                         context_impl &Context,
                          kernel_bundle_impl *KernelBundleImpl,
                          const KernelArgMask *ArgMask)
-    : MKernel(Kernel), MContext(Context.shared_from_this()),
-      MProgram(ProgramManager::getInstance().getUrProgramFromUrKernel(Kernel,
+    : MKernel(std::move(Kernel)), MContext(Context.shared_from_this()),
+      MProgram(ProgramManager::getInstance().getUrProgramFromUrKernel(MKernel,
                                                                       Context)),
       MCreatedFromSource(true),
       MKernelBundleImpl(KernelBundleImpl ? KernelBundleImpl->shared_from_this()
@@ -39,12 +40,13 @@ kernel_impl::kernel_impl(ur_kernel_handle_t Kernel, context_impl &Context,
   enableUSMIndirectAccess();
 }
 
-kernel_impl::kernel_impl(ur_kernel_handle_t Kernel, context_impl &ContextImpl,
+kernel_impl::kernel_impl(Managed<ur_kernel_handle_t> &&Kernel,
+                         context_impl &ContextImpl,
                          std::shared_ptr<device_image_impl> &&DeviceImageImpl,
                          const kernel_bundle_impl &KernelBundleImpl,
                          const KernelArgMask *ArgMask,
                          ur_program_handle_t Program, std::mutex *CacheMutex)
-    : MKernel(Kernel), MContext(ContextImpl.shared_from_this()),
+    : MKernel(std::move(Kernel)), MContext(ContextImpl.shared_from_this()),
       MProgram(Program),
       MCreatedFromSource(DeviceImageImpl->isNonSYCLSourceBased()),
       MDeviceImageImpl(std::move(DeviceImageImpl)),
@@ -58,14 +60,21 @@ kernel_impl::kernel_impl(ur_kernel_handle_t Kernel, context_impl &ContextImpl,
     enableUSMIndirectAccess();
 }
 
-kernel_impl::~kernel_impl() {
-  try {
-    // TODO catch an exception and put it to list of asynchronous exceptions
-    getAdapter().call<UrApiKind::urKernelRelease>(MKernel);
-  } catch (std::exception &e) {
-    __SYCL_REPORT_EXCEPTION_TO_STREAM("exception in ~kernel_impl", e);
-  }
+#ifdef _MSC_VER
+#pragma warning(push)
+// https://developercommunity.visualstudio.com/t/False-C4297-warning-while-using-function/1130300
+// https://godbolt.org/z/xsMvKf84f
+#pragma warning(disable : 4297)
+#endif
+kernel_impl::~kernel_impl() try {
+} catch (std::exception &e) {
+  // TODO put it to list of asynchronous exceptions
+  __SYCL_REPORT_EXCEPTION_TO_STREAM("exception in ~kernel_impl", e);
+  return; // Don't re-throw.
 }
+#ifdef _MSC_VER
+#pragma warning(pop)
+#endif
 
 bool kernel_impl::isCreatedFromSource() const {
   // TODO it is not clear how to understand whether the SYCL kernel is created
diff --git a/sycl/source/detail/kernel_impl.hpp b/sycl/source/detail/kernel_impl.hpp
@@ -39,7 +39,7 @@ class kernel_impl {
   /// \param Kernel is a valid UrKernel instance
   /// \param Context is a valid SYCL context
   /// \param KernelBundleImpl is a valid instance of kernel_bundle_impl
-  kernel_impl(ur_kernel_handle_t Kernel, context_impl &Context,
+  kernel_impl(Managed<ur_kernel_handle_t> &&Kernel, context_impl &Context,
               kernel_bundle_impl *KernelBundleImpl,
               const KernelArgMask *ArgMask = nullptr);
 
@@ -49,7 +49,7 @@ class kernel_impl {
   /// \param Kernel is a valid UrKernel instance
   /// \param ContextImpl is a valid SYCL context
   /// \param KernelBundleImpl is a valid instance of kernel_bundle_impl
-  kernel_impl(ur_kernel_handle_t Kernel, context_impl &ContextImpl,
+  kernel_impl(Managed<ur_kernel_handle_t> &&Kernel, context_impl &ContextImpl,
               std::shared_ptr<device_image_impl> &&DeviceImageImpl,
               const kernel_bundle_impl &KernelBundleImpl,
               const KernelArgMask *ArgMask, ur_program_handle_t Program,
@@ -198,11 +198,7 @@ class kernel_impl {
   typename Param::return_type ext_oneapi_get_info(queue Queue,
                                                   const range<1> &WG) const;
 
-  /// Get a constant reference to a raw kernel object.
-  ///
-  /// \return a constant reference to a valid UrKernel instance with raw
-  /// kernel object.
-  const ur_kernel_handle_t &getHandleRef() const { return MKernel; }
+  ur_kernel_handle_t getHandleRef() const { return MKernel; }
 
   /// Check if kernel was created from a program that had been created from
   /// source.
@@ -243,7 +239,7 @@ class kernel_impl {
   std::string_view getName() const;
 
 private:
-  ur_kernel_handle_t MKernel = nullptr;
+  Managed<ur_kernel_handle_t> MKernel;
   const std::shared_ptr<context_impl> MContext;
   const ur_program_handle_t MProgram = nullptr;
   bool MCreatedFromSource = true;
diff --git a/sycl/source/detail/kernel_name_based_cache_t.hpp b/sycl/source/detail/kernel_name_based_cache_t.hpp
@@ -22,7 +22,7 @@ namespace detail {
 using FastKernelCacheKeyT = std::pair<ur_device_handle_t, ur_context_handle_t>;
 
 struct FastKernelCacheVal {
-  ur_kernel_handle_t MKernelHandle;    /* UR kernel handle pointer. */
+  Managed<ur_kernel_handle_t> MKernelHandle; /* UR kernel. */
   std::mutex *MMutex;                  /* Mutex guarding this kernel. When
                                      caching is disabled, the pointer is
                                      nullptr. */
@@ -34,18 +34,15 @@ struct FastKernelCacheVal {
                             cache is destroyed deliberately before the
                             adapter. */
 
-  FastKernelCacheVal(ur_kernel_handle_t KernelHandle, std::mutex *Mutex,
-                     const KernelArgMask *KernelArgMask,
+  FastKernelCacheVal(Managed<ur_kernel_handle_t> &&KernelHandle,
+                     std::mutex *Mutex, const KernelArgMask *KernelArgMask,
                      Managed<ur_program_handle_t> &&ProgramHandle,
                      adapter_impl &Adapter)
-      : MKernelHandle(KernelHandle), MMutex(Mutex),
+      : MKernelHandle(std::move(KernelHandle)), MMutex(Mutex),
         MKernelArgMask(KernelArgMask), MProgramHandle(std::move(ProgramHandle)),
         MAdapter(Adapter) {}
 
   ~FastKernelCacheVal() {
-    if (MKernelHandle)
-      MAdapter.call<sycl::detail::UrApiKind::urKernelRelease>(MKernelHandle);
-    MKernelHandle = nullptr;
     MMutex = nullptr;
     MKernelArgMask = nullptr;
   }
diff --git a/sycl/source/detail/kernel_program_cache.hpp b/sycl/source/detail/kernel_program_cache.hpp
@@ -194,22 +194,23 @@ class KernelProgramCache {
 
   struct KernelBuildResult
       : public BuildResult<
-            std::pair<ur_kernel_handle_t, const KernelArgMask *>> {
-    const adapter_impl &MAdapter;
-    KernelBuildResult(const adapter_impl &Adapter) : MAdapter(Adapter) {
-      Val.first = nullptr;
-    }
-    ~KernelBuildResult() {
-      try {
-        if (Val.first) {
-          ur_result_t Err =
-              MAdapter.call_nocheck<UrApiKind::urKernelRelease>(Val.first);
-          __SYCL_CHECK_UR_CODE_NO_EXC(Err, MAdapter.getBackend());
-        }
-      } catch (std::exception &e) {
-        __SYCL_REPORT_EXCEPTION_TO_STREAM("exception in ~KernelBuildResult", e);
-      }
+            std::pair<Managed<ur_kernel_handle_t>, const KernelArgMask *>> {
+    KernelBuildResult() = default;
+
+#ifdef _MSC_VER
+#pragma warning(push)
+// https://developercommunity.visualstudio.com/t/False-C4297-warning-while-using-function/1130300
+// https://godbolt.org/z/xsMvKf84f
+#pragma warning(disable : 4297)
+#endif
+    ~KernelBuildResult() try {
+    } catch (std::exception &e) {
+      __SYCL_REPORT_EXCEPTION_TO_STREAM("exception in ~KernelBuildResult", e);
+      return; // Don't re-throw.
     }
+#ifdef _MSC_VER
+#pragma warning(pop)
+#endif
   };
 
   using KernelCacheT = emhash8::HashMap<
@@ -445,7 +446,7 @@ class KernelProgramCache {
     auto &Cache = LockedCache.get()[Program];
     auto [It, DidInsert] = Cache.try_emplace(KernelName, nullptr);
     if (DidInsert) {
-      It->second = std::make_shared<KernelBuildResult>(getAdapter());
+      It->second = std::make_shared<KernelBuildResult>();
       traceKernel("Kernel inserted.", KernelName);
     } else
       traceKernel("Kernel fetched.", KernelName);
diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp
@@ -1107,9 +1107,8 @@ FastKernelCacheValPtr ProgramManager::getOrCreateKernel(
       getBuiltURProgram(ContextImpl, DeviceImpl, KernelName, NDRDesc);
 
   auto BuildF = [this, &Program, &KernelName, &ContextImpl] {
-    ur_kernel_handle_t Kernel = nullptr;
-
     adapter_impl &Adapter = ContextImpl.getAdapter();
+    Managed<ur_kernel_handle_t> Kernel{Adapter};
     Adapter.call<errc::kernel_not_supported, UrApiKind::urKernelCreate>(
         Program, KernelName.data(), &Kernel);
 
@@ -1126,7 +1125,7 @@ FastKernelCacheValPtr ProgramManager::getOrCreateKernel(
     const KernelArgMask *ArgMask = nullptr;
     if (!m_UseSpvFile)
       ArgMask = getEliminatedKernelArgMask(Program, KernelName);
-    return std::make_pair(Kernel, ArgMask);
+    return std::make_pair(std::move(Kernel), ArgMask);
   };
 
   auto GetCachedBuildF = [&Cache, &KernelName, &Program]() {
@@ -1138,24 +1137,19 @@ FastKernelCacheValPtr ProgramManager::getOrCreateKernel(
     // threads when caching is disabled, so we can return
     // nullptr for the mutex.
     auto [Kernel, ArgMask] = BuildF();
-    return std::make_shared<FastKernelCacheVal>(
-        Kernel, nullptr, ArgMask, std::move(Program), ContextImpl.getAdapter());
+    return std::make_shared<FastKernelCacheVal>(std::move(Kernel), nullptr,
+                                                ArgMask, std::move(Program),
+                                                ContextImpl.getAdapter());
   }
 
   std::shared_ptr<KernelProgramCache::KernelBuildResult> BuildResult =
       Cache.getOrBuild<errc::invalid>(GetCachedBuildF, BuildF);
   assert(BuildResult && "getOrBuild isn't supposed to return nullptr!");
-  const std::pair<ur_kernel_handle_t, const KernelArgMask *>
+  std::pair<Managed<ur_kernel_handle_t>, const KernelArgMask *>
       &KernelArgMaskPair = BuildResult->Val;
   auto ret_val = std::make_shared<FastKernelCacheVal>(
-      KernelArgMaskPair.first, &(BuildResult->MBuildResultMutex),
+      KernelArgMaskPair.first.retain(), &(BuildResult->MBuildResultMutex),
       KernelArgMaskPair.second, std::move(Program), ContextImpl.getAdapter());
-  // If caching is enabled, one copy of the kernel handle will be
-  // stored in FastKernelCacheVal, and one is in
-  // KernelProgramCache::MKernelsPerProgramCache. To cover
-  // MKernelsPerProgramCache, we need to increase the ref count of the kernel.
-  ContextImpl.getAdapter().call<UrApiKind::urKernelRetain>(
-      KernelArgMaskPair.first);
   Cache.saveKernel(KernelName, UrDevice, ret_val, CacheHintPtr);
   return ret_val;
 }
@@ -3116,7 +3110,7 @@ ProgramManager::build(const DevImgPlainWithDeps &DevImgWithDeps,
 
 // When caching is enabled, the returned UrKernel will already have
 // its ref count incremented.
-std::tuple<ur_kernel_handle_t, std::mutex *, const KernelArgMask *>
+std::tuple<Managed<ur_kernel_handle_t>, std::mutex *, const KernelArgMask *>
 ProgramManager::getOrCreateKernel(const context &Context,
                                   KernelNameStrRefT KernelName,
                                   const property_list &PropList,
@@ -3133,9 +3127,9 @@ ProgramManager::getOrCreateKernel(const context &Context,
   KernelProgramCache &Cache = Ctx.getKernelProgramCache();
 
   auto BuildF = [this, &Program, &KernelName, &Ctx] {
-    ur_kernel_handle_t Kernel = nullptr;
-
     adapter_impl &Adapter = Ctx.getAdapter();
+    Managed<ur_kernel_handle_t> Kernel{Adapter};
+
     Adapter.call<UrApiKind::urKernelCreate>(Program, KernelName.data(),
                                             &Kernel);
 
@@ -3152,7 +3146,7 @@ ProgramManager::getOrCreateKernel(const context &Context,
     const KernelArgMask *KernelArgMask =
         getEliminatedKernelArgMask(Program, KernelName);
 
-    return std::make_pair(Kernel, KernelArgMask);
+    return std::make_pair(std::move(Kernel), KernelArgMask);
   };
 
   auto GetCachedBuildF = [&Cache, &KernelName, Program]() {
@@ -3164,7 +3158,7 @@ ProgramManager::getOrCreateKernel(const context &Context,
     // threads when caching is disabled, so we can return
     // nullptr for the mutex.
     auto [Kernel, ArgMask] = BuildF();
-    return make_tuple(Kernel, nullptr, ArgMask);
+    return make_tuple(std::move(Kernel), nullptr, ArgMask);
   }
 
   std::shared_ptr<KernelProgramCache::KernelBuildResult> BuildResult =
@@ -3174,8 +3168,7 @@ ProgramManager::getOrCreateKernel(const context &Context,
   // stored in the cache, and one handle is returned to the
   // caller. In that case, we need to increase the ref count of the
   // kernel.
-  Ctx.getAdapter().call<UrApiKind::urKernelRetain>(BuildResult->Val.first);
-  return std::make_tuple(BuildResult->Val.first,
+  return std::make_tuple(BuildResult->Val.first.retain(),
                          &(BuildResult->MBuildResultMutex),
                          BuildResult->Val.second);
 }
@@ -3242,15 +3235,17 @@ ur_kernel_handle_t ProgramManager::getOrCreateMaterializedKernel(
       build(std::move(ProgramManaged), ContextImpl, CompileOpts, LinkOpts, Devs,
             /*For non SPIR-V devices DeviceLibReqdMask is always 0*/ 0,
             ExtraProgramsToLink);
-  ur_kernel_handle_t UrKernel{nullptr};
+  Managed<ur_kernel_handle_t> UrKernel{Adapter};
   Adapter.call<errc::kernel_not_supported, UrApiKind::urKernelCreate>(
       BuildProgram, KernelName.data(), &UrKernel);
+  ur_kernel_handle_t RawUrKernel = UrKernel;
   {
     std::lock_guard<std::mutex> KernelIDsGuard(m_KernelIDsMutex);
-    m_MaterializedKernels[KernelName][SpecializationConsts] = UrKernel;
+    m_MaterializedKernels[KernelName][SpecializationConsts] =
+        std::move(UrKernel);
   }
 
-  return UrKernel;
+  return RawUrKernel;
 }
 
 bool doesDevSupportDeviceRequirements(const device_impl &Dev,
diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp
diff --git a/sycl/source/kernel.cpp b/sycl/source/kernel.cpp
diff --git a/sycl/test-e2e/KernelAndProgram/disable-caching.cpp b/sycl/test-e2e/KernelAndProgram/disable-caching.cpp
