From 2e777b50460f598e7d37b09d36c78bf171d786a0 Mon Sep 17 00:00:00 2001
From: Steffen Larsen <steffen.larsen@intel.com>
Date: Thu, 5 Sep 2024 17:14:14 +0200
Subject: [PATCH] [SYCL] Add host kernel instantiation for debuggers (#15256)

The changes in https://github.com/intel/llvm/pull/14460 removed the
seemingly unused functions for running kernels on on host. However, this
turned out to be used by debuggers as they need the kernel code to be in
the host executable.

This commit adds a simplified version of the kernel instantiation that
the aforementioned patch removed.

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 sycl/include/sycl/detail/cg_types.hpp | 46 +++++++++++++++++++++++++++
 sycl/test/abi/vtable.cpp              | 12 +++++++
 2 files changed, 58 insertions(+)

diff --git a/sycl/include/sycl/detail/cg_types.hpp b/sycl/include/sycl/detail/cg_types.hpp
index 7db3171ba28fc..ef0ef198c40d0 100644
--- a/sycl/include/sycl/detail/cg_types.hpp
+++ b/sycl/include/sycl/detail/cg_types.hpp
@@ -154,6 +154,8 @@ runKernelWithArg(KernelType KernelName, ArgType Arg) {
 // The pure virtual class aimed to store lambda/functors of any type.
 class HostKernelBase {
 public:
+  // NOTE: InstatitateKernelOnHost() should not be called.
+  virtual void InstatitateKernelOnHost() = 0;
   // Return pointer to the lambda object.
   // Used to extract captured variables.
   virtual char *getPtr() = 0;
@@ -172,6 +174,50 @@ class HostKernel : public HostKernelBase {
 public:
   HostKernel(KernelType Kernel) : MKernel(Kernel) {}
 
+  // 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: InstatitateKernelOnHost() should not be called.
+  void InstatitateKernelOnHost() override {
+    if constexpr (std::is_same_v<KernelArgType, void>) {
+      runKernelWithoutArg(MKernel);
+    } 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);
+    } 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>>;
+      KernelArgType Item = IDBuilder::createItem<Dims, HasOffset>(
+          InitializedVal<Dims, range>::template get<1>(),
+          InitializedVal<Dims, id>::template get<0>());
+      runKernelWithArg<KernelArgType>(MKernel, Item);
+    } 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);
+    } 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);
+    } 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{});
+    }
+  }
+
   char *getPtr() override { return reinterpret_cast<char *>(&MKernel); }
 
   ~HostKernel() = default;
diff --git a/sycl/test/abi/vtable.cpp b/sycl/test/abi/vtable.cpp
index f1ff7e02f1900..b3d16e3c4381d 100644
--- a/sycl/test/abi/vtable.cpp
+++ b/sycl/test/abi/vtable.cpp
@@ -8,6 +8,18 @@
 // Changing vtable breaks ABI. If this test fails, please, refer to ABI Policy
 // Guide for further instructions.
 
+void foo(sycl::detail::HostKernelBase &HKB) {
+  HKB.InstatitateKernelOnHost();
+}
+// CHECK:      Vtable for 'sycl::detail::HostKernelBase' (6 entries).
+// CHECK-NEXT:   0 | offset_to_top (0)
+// CHECK-NEXT:   1 | sycl::detail::HostKernelBase RTTI
+// CHECK-NEXT:       -- (sycl::detail::HostKernelBase, 0) vtable address --
+// CHECK-NEXT:   2 | void sycl::detail::HostKernelBase::InstatitateKernelOnHost() [pure]
+// CHECK-NEXT:   3 | char *sycl::detail::HostKernelBase::getPtr() [pure]
+// CHECK-NEXT:   4 | sycl::detail::HostKernelBase::~HostKernelBase() [complete]
+// CHECK-NEXT:   5 | sycl::detail::HostKernelBase::~HostKernelBase() [deleting]
+
 void foo(sycl::detail::PropertyWithDataBase *Prop) { delete Prop; }
 // CHECK:    Vtable for 'sycl::detail::PropertyWithDataBase' (4 entries).
 // CHECK-NEXT:   0 | offset_to_top (0)