[SYCL] Postpone creation of HostKernel copy

Alexandr-Konovalov · Alexandr-Konovalov · commit cc71f38dac91 · 2025-09-29T18:46:19.000+02:00
Create possibility to delay creation of copy of HostKernel till it became used
out of submit stack, i.e. by scheduler. Do type erasure for kernel lambda via
vptr in HostKernelRefBase.
diff --git a/sycl/include/sycl/detail/cg_types.hpp b/sycl/include/sycl/detail/cg_types.hpp
@@ -235,6 +235,93 @@ class HostKernel : public HostKernelBase {
 #endif
 };
 
+// the class keeps reference to a lambda allocated externally on stack
+class HostKernelRefBase : public HostKernelBase {
+public:
+  virtual std::shared_ptr<HostKernelBase> takeOrCopyOwnership() const = 0;
+};
+
+template <class KernelType, class KernelArgType, int Dims>
+class HostKernelRef : public HostKernelRefBase {
+  const KernelType &MKernel;
+
+public:
+  HostKernelRef(const KernelType &Kernel) : MKernel(Kernel) {}
+
+  virtual char *getPtr() override {
+    return const_cast<char *>(reinterpret_cast<const char *>(&MKernel));
+  }
+  virtual std::shared_ptr<HostKernelBase> takeOrCopyOwnership() const override {
+    std::shared_ptr<HostKernelBase> Kernel;
+    Kernel.reset(new HostKernel<KernelType, KernelArgType, Dims>(MKernel));
+    return Kernel;
+  }
+
+  ~HostKernelRef() noexcept override = default;
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+  // This function is needed for host-side compilation to keep kernels
+  // instantitated. This is important for debuggers to be able to associate
+  // kernel code instructions with source code lines.
+  // NOTE: InstatiateKernelOnHost() should not be called.
+  void InstantiateKernelOnHost() override {
+    using IDBuilder = sycl::detail::Builder;
+    constexpr bool HasKernelHandlerArg =
+        KernelLambdaHasKernelHandlerArgT<KernelType, KernelArgType>::value;
+    if constexpr (std::is_same_v<KernelArgType, void>) {
+      runKernelWithoutArg(MKernel, std::bool_constant<HasKernelHandlerArg>());
+    } else if constexpr (std::is_same_v<KernelArgType, sycl::id<Dims>>) {
+      sycl::id ID = InitializedVal<Dims, id>::template get<0>();
+      runKernelWithArg<const KernelArgType &>(
+          MKernel, ID, std::bool_constant<HasKernelHandlerArg>());
+    } else if constexpr (std::is_same_v<KernelArgType, item<Dims, true>> ||
+                         std::is_same_v<KernelArgType, item<Dims, false>>) {
+      constexpr bool HasOffset =
+          std::is_same_v<KernelArgType, item<Dims, true>>;
+      if constexpr (!HasOffset) {
+        KernelArgType Item = IDBuilder::createItem<Dims, HasOffset>(
+            InitializedVal<Dims, range>::template get<1>(),
+            InitializedVal<Dims, id>::template get<0>());
+        runKernelWithArg<KernelArgType>(
+            MKernel, Item, std::bool_constant<HasKernelHandlerArg>());
+      } else {
+        KernelArgType Item = IDBuilder::createItem<Dims, HasOffset>(
+            InitializedVal<Dims, range>::template get<1>(),
+            InitializedVal<Dims, id>::template get<0>(),
+            InitializedVal<Dims, id>::template get<0>());
+        runKernelWithArg<KernelArgType>(
+            MKernel, Item, std::bool_constant<HasKernelHandlerArg>());
+      }
+    } else if constexpr (std::is_same_v<KernelArgType, nd_item<Dims>>) {
+      sycl::range<Dims> Range = InitializedVal<Dims, range>::template get<1>();
+      sycl::id<Dims> ID = InitializedVal<Dims, id>::template get<0>();
+      sycl::group<Dims> Group =
+          IDBuilder::createGroup<Dims>(Range, Range, Range, ID);
+      sycl::item<Dims, true> GlobalItem =
+          IDBuilder::createItem<Dims, true>(Range, ID, ID);
+      sycl::item<Dims, false> LocalItem =
+          IDBuilder::createItem<Dims, false>(Range, ID);
+      KernelArgType NDItem =
+          IDBuilder::createNDItem<Dims>(GlobalItem, LocalItem, Group);
+      runKernelWithArg<const KernelArgType>(
+          MKernel, NDItem, std::bool_constant<HasKernelHandlerArg>());
+    } else if constexpr (std::is_same_v<KernelArgType, sycl::group<Dims>>) {
+      sycl::range<Dims> Range = InitializedVal<Dims, range>::template get<1>();
+      sycl::id<Dims> ID = InitializedVal<Dims, id>::template get<0>();
+      KernelArgType Group =
+          IDBuilder::createGroup<Dims>(Range, Range, Range, ID);
+      runKernelWithArg<KernelArgType>(
+          MKernel, Group, std::bool_constant<HasKernelHandlerArg>());
+    } else {
+      // Assume that anything else can be default-constructed. If not, this
+      // should fail to compile and the implementor should implement a generic
+      // case for the new argument type.
+      runKernelWithArg<KernelArgType>(
+          MKernel, KernelArgType{}, std::bool_constant<HasKernelHandlerArg>());
+    }
+  }
+#endif
+};
+
 // This function is needed for host-side compilation to keep kernels
 // instantitated. This is important for debuggers to be able to associate
 // kernel code instructions with source code lines.
diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp
@@ -65,14 +65,14 @@ auto get_native(const SyclObjectT &Obj)
 template <int Dims>
 event __SYCL_EXPORT submit_kernel_direct_with_event_impl(
     const queue &Queue, const nd_range<Dims> &Range,
-    std::shared_ptr<detail::HostKernelBase> &HostKernel,
+    detail::HostKernelRefBase &HostKernel,
     detail::DeviceKernelInfo *DeviceKernelInfo,
     const detail::code_location &CodeLoc, bool IsTopCodeLoc);
 
 template <int Dims>
 void __SYCL_EXPORT submit_kernel_direct_without_event_impl(
     const queue &Queue, const nd_range<Dims> &Range,
-    std::shared_ptr<detail::HostKernelBase> &HostKernel,
+    detail::HostKernelRefBase &HostKernel,
     detail::DeviceKernelInfo *DeviceKernelInfo,
     const detail::code_location &CodeLoc, bool IsTopCodeLoc);
 
@@ -180,8 +180,15 @@ auto submit_kernel_direct(
       "must be either sycl::nd_item or be convertible from sycl::nd_item");
   using TransformedArgType = sycl::nd_item<Dims>;
 
-  std::shared_ptr<detail::HostKernelBase> HostKernel = std::make_shared<
-      detail::HostKernel<KernelType, TransformedArgType, Dims>>(KernelFunc);
+  HostKernelRef<KernelType, TransformedArgType, Dims> HostKernel(KernelFunc);
+
+  // Instantiating the kernel on the host improves debugging.
+  // Passing this pointer to another translation unit prevents optimization.
+#ifndef NDEBUG
+  // TODO: call library to prevent dropping call due to optimization
+  (void)detail::GetInstantiateKernelOnHostPtr<KernelType, LambdaArgType,
+                                              Dims>();
+#endif
 
   detail::DeviceKernelInfo *DeviceKernelInfoPtr =
       &detail::getDeviceKernelInfo<NameT>();
diff --git a/sycl/source/detail/queue_impl.cpp b/sycl/source/detail/queue_impl.cpp
@@ -422,14 +422,16 @@ queue_impl::submit_impl(const detail::type_erased_cgfo_ty &CGF,
 
 detail::EventImplPtr queue_impl::submit_kernel_direct_impl(
     const NDRDescT &NDRDesc,
-    std::shared_ptr<detail::HostKernelBase> &HostKernel,
+    detail::HostKernelRefBase &HostKernel,
     detail::DeviceKernelInfo *DeviceKernelInfo, bool CallerNeedsEvent,
     const detail::code_location &CodeLoc, bool IsTopCodeLoc) {
 
   KernelData KData;
 
+  std::shared_ptr<detail::HostKernelBase> HostKernelPtr = HostKernel.takeOrCopyOwnership();
+
   KData.setDeviceKernelInfoPtr(DeviceKernelInfo);
-  KData.setKernelFunc(HostKernel->getPtr());
+  KData.setKernelFunc(HostKernelPtr->getPtr());
   KData.setNDRDesc(NDRDesc);
 
   auto SubmitKernelFunc =
@@ -441,7 +443,7 @@ detail::EventImplPtr queue_impl::submit_kernel_direct_impl(
     KData.extractArgsAndReqsFromLambda();
 
     CommandGroup.reset(new detail::CGExecKernel(
-        KData.getNDRDesc(), HostKernel,
+        KData.getNDRDesc(), std::move(HostKernelPtr),
         nullptr, // Kernel
         nullptr, // KernelBundle
         std::move(CGData), std::move(KData).getArgs(),
diff --git a/sycl/source/detail/queue_impl.hpp b/sycl/source/detail/queue_impl.hpp
@@ -362,7 +362,7 @@ class queue_impl : public std::enable_shared_from_this<queue_impl> {
   template <int Dims>
   event submit_kernel_direct_with_event(
       const nd_range<Dims> &Range,
-      std::shared_ptr<detail::HostKernelBase> &HostKernel,
+      detail::HostKernelRefBase &HostKernel,
       detail::DeviceKernelInfo *DeviceKernelInfo,
       const detail::code_location &CodeLoc, bool IsTopCodeLoc) {
     detail::EventImplPtr EventImpl =
@@ -374,7 +374,7 @@ class queue_impl : public std::enable_shared_from_this<queue_impl> {
   template <int Dims>
   void submit_kernel_direct_without_event(
       const nd_range<Dims> &Range,
-      std::shared_ptr<detail::HostKernelBase> &HostKernel,
+      detail::HostKernelRefBase &HostKernel,
       detail::DeviceKernelInfo *DeviceKernelInfo,
       const detail::code_location &CodeLoc, bool IsTopCodeLoc) {
     submit_kernel_direct_impl(NDRDescT{Range}, HostKernel, DeviceKernelInfo,
@@ -906,7 +906,7 @@ class queue_impl : public std::enable_shared_from_this<queue_impl> {
   /// \return a SYCL event representing submitted command group or nullptr.
   detail::EventImplPtr submit_kernel_direct_impl(
       const NDRDescT &NDRDesc,
-      std::shared_ptr<detail::HostKernelBase> &HostKernel,
+      detail::HostKernelRefBase &HostKernel,
       detail::DeviceKernelInfo *DeviceKernelInfo, bool CallerNeedsEvent,
       const detail::code_location &CodeLoc, bool IsTopCodeLoc);
 
diff --git a/sycl/source/queue.cpp b/sycl/source/queue.cpp
@@ -466,7 +466,7 @@ const property_list &queue::getPropList() const { return impl->getPropList(); }
 template <int Dims>
 event submit_kernel_direct_with_event_impl(
     const queue &Queue, const nd_range<Dims> &Range,
-    std::shared_ptr<detail::HostKernelBase> &HostKernel,
+    detail::HostKernelRefBase &HostKernel,
     detail::DeviceKernelInfo *DeviceKernelInfo,
     const detail::code_location &CodeLoc, bool IsTopCodeLoc) {
   return getSyclObjImpl(Queue)->submit_kernel_direct_with_event(
@@ -475,26 +475,26 @@ event submit_kernel_direct_with_event_impl(
 
 template event __SYCL_EXPORT submit_kernel_direct_with_event_impl<1>(
     const queue &Queue, const nd_range<1> &Range,
-    std::shared_ptr<detail::HostKernelBase> &HostKernel,
+    detail::HostKernelRefBase &HostKernel,
     detail::DeviceKernelInfo *DeviceKernelInfo,
     const detail::code_location &CodeLoc, bool IsTopCodeLoc);
 
 template event __SYCL_EXPORT submit_kernel_direct_with_event_impl<2>(
     const queue &Queue, const nd_range<2> &Range,
-    std::shared_ptr<detail::HostKernelBase> &HostKernel,
+    detail::HostKernelRefBase &HostKernel,
     detail::DeviceKernelInfo *DeviceKernelInfo,
     const detail::code_location &CodeLoc, bool IsTopCodeLoc);
 
 template event __SYCL_EXPORT submit_kernel_direct_with_event_impl<3>(
     const queue &Queue, const nd_range<3> &Range,
-    std::shared_ptr<detail::HostKernelBase> &HostKernel,
+    detail::HostKernelRefBase &HostKernel,
     detail::DeviceKernelInfo *DeviceKernelInfo,
     const detail::code_location &CodeLoc, bool IsTopCodeLoc);
 
 template <int Dims>
 void submit_kernel_direct_without_event_impl(
     const queue &Queue, const nd_range<Dims> &Range,
-    std::shared_ptr<detail::HostKernelBase> &HostKernel,
+    detail::HostKernelRefBase &HostKernel,
     detail::DeviceKernelInfo *DeviceKernelInfo,
     const detail::code_location &CodeLoc, bool IsTopCodeLoc) {
   getSyclObjImpl(Queue)->submit_kernel_direct_without_event(
@@ -503,19 +503,19 @@ void submit_kernel_direct_without_event_impl(
 
 template void __SYCL_EXPORT submit_kernel_direct_without_event_impl<1>(
     const queue &Queue, const nd_range<1> &Range,
-    std::shared_ptr<detail::HostKernelBase> &HostKernel,
+    detail::HostKernelRefBase &HostKernel,
     detail::DeviceKernelInfo *DeviceKernelInfo,
     const detail::code_location &CodeLoc, bool IsTopCodeLoc);
 
 template void __SYCL_EXPORT submit_kernel_direct_without_event_impl<2>(
     const queue &Queue, const nd_range<2> &Range,
-    std::shared_ptr<detail::HostKernelBase> &HostKernel,
+    detail::HostKernelRefBase &HostKernel,
     detail::DeviceKernelInfo *DeviceKernelInfo,
     const detail::code_location &CodeLoc, bool IsTopCodeLoc);
 
 template void __SYCL_EXPORT submit_kernel_direct_without_event_impl<3>(
     const queue &Queue, const nd_range<3> &Range,
-    std::shared_ptr<detail::HostKernelBase> &HostKernel,
+    detail::HostKernelRefBase &HostKernel,
     detail::DeviceKernelInfo *DeviceKernelInfo,
     const detail::code_location &CodeLoc, bool IsTopCodeLoc);
 
