[SYCL] Fix unintentional vtable ABI break in HostKernelBase (#15310)

The changes in #15256 unintentionally
broke ABI by reordering the virtual member functions of HostKernelBase.
This commit amends this by moving the new function to the end of the new
member function.

Additionally, this fixes a typo in the new function.

---------

Signed-off-by: Larsen, Steffen <[email protected]>
Co-authored-by: Marcos Maronas <[email protected]>

Author: steffenlarsen

sycl/include/sycl/detail/cg_types.hpp

@@ -154,12 +154,12 @@ 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;
   virtual ~HostKernelBase() = default;
+  // NOTE: InstatiateKernelOnHost() should not be called.
+  virtual void InstantiateKernelOnHost() = 0;
 };
 
 // Class which stores specific lambda object.
@@ -174,11 +174,15 @@ class HostKernel : public HostKernelBase {
 public:
   HostKernel(KernelType Kernel) : MKernel(Kernel) {}
 
+  char *getPtr() override { return reinterpret_cast<char *>(&MKernel); }
+
+  ~HostKernel() = default;
+
   // 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 {
+  // NOTE: InstatiateKernelOnHost() should not be called.
+  void InstantiateKernelOnHost() override {
     if constexpr (std::is_same_v<KernelArgType, void>) {
       runKernelWithoutArg(MKernel);
     } else if constexpr (std::is_same_v<KernelArgType, sycl::id<Dims>>) {
@@ -217,10 +221,6 @@ class HostKernel : public HostKernelBase {
       runKernelWithArg<KernelArgType>(MKernel, KernelArgType{});
     }
   }
-
-  char *getPtr() override { return reinterpret_cast<char *>(&MKernel); }
-
-  ~HostKernel() = default;
 };
 
 } // namespace detail

sycl/test/abi/vtable.cpp

@@ -9,16 +9,16 @@
 // Guide for further instructions.
 
 void foo(sycl::detail::HostKernelBase &HKB) {
-  HKB.InstatitateKernelOnHost();
+  HKB.InstantiateKernelOnHost();
 }
 // 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]
+// CHECK-NEXT:   2 | char *sycl::detail::HostKernelBase::getPtr() [pure]
+// CHECK-NEXT:   3 | sycl::detail::HostKernelBase::~HostKernelBase() [complete]
+// CHECK-NEXT:   4 | sycl::detail::HostKernelBase::~HostKernelBase() [deleting]
+// CHECK-NEXT:   5 | void sycl::detail::HostKernelBase::InstantiateKernelOnHost() [pure]
 
 void foo(sycl::detail::PropertyWithDataBase *Prop) { delete Prop; }
 // CHECK:    Vtable for 'sycl::detail::PropertyWithDataBase' (4 entries).