[SYCL] Add host kernel instantiation for debuggers (#15256)

The changes in #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

@@ -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;

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)