[SYCL][NFCI] Drop ESIMD emulator leftovers

#4020 introduced some kind of wrapper for SYCL kernels on host
to be able to launch them on ESIMD emulator backend. Even though we had
dropped that feature since then, we didn't remove corresponding
`handler.hpp` modifications. This PR removes them.

The main effect expected from this PR is compilation time improvement:
those kernel wrappers are specialized by kernel name, meaning that there
will be plenty of extra useless functions emitted during host compilation
pass for every kernel in a program.

This patch also uncovered a missing case in
`HostKernel::InstantiateKernelOnHost` which we couldn't ever encounter
because we transformed a host kernel to always accept `nd_item`, thus
always skipping problematic `item` code path.

This PR is not expected to introduce any functional changes, but since
I'm not very familiar with the SYCL RT, I'm not entirely sure of that.

Author: AlexeySachkov

sycl/include/sycl/detail/cg_types.hpp

@@ -192,10 +192,18 @@ class HostKernel : public HostKernelBase {
                          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);
+      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);
+      } 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);
+      }
     } 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/include/sycl/detail/helpers.hpp

@@ -10,8 +10,8 @@
 
 #include <CL/__spirv/spirv_types.hpp> // for MemorySemanticsMask
 #include <sycl/access/access.hpp>     // for fence_space
-#include <sycl/detail/export.hpp>     // for __SYCL_EXPORT
-#include <sycl/memory_enums.hpp>      // for memory_order
+#include <sycl/detail/export.hpp> // for __SYCL_EXPORT
+#include <sycl/memory_enums.hpp>  // for memory_order
 
 #ifdef __SYCL_DEVICE_ONLY__
 #include <CL/__spirv/spirv_vars.hpp>

sycl/include/sycl/handler.hpp

@@ -767,130 +767,6 @@ class __SYCL_EXPORT handler {
           &DynamicParamBase,
       int ArgIndex);
 
-  /* The kernel passed to StoreLambda can take an id, an item or an nd_item as
-   * its argument. Since esimd adapter directly invokes the kernel (doesn’t use
-   * urKernelSetArg), the kernel argument type must be known to the adapter.
-   * However, passing kernel argument type to the adapter requires changing ABI
-   * in HostKernel class. To overcome this problem, helpers below wrap the
-   * “original” kernel with a functor that always takes an nd_item as argument.
-   * A functor is used instead of a lambda because extractArgsAndReqsFromLambda
-   * needs access to the “original” kernel and keeps references to its internal
-   * data, i.e. the kernel passed as argument cannot be local in scope. The
-   * functor itself is again encapsulated in a std::function since functor’s
-   * type is unknown to the adapter.
-   */
-
-  // For 'id, item w/wo offset, nd_item' kernel arguments
-  template <class KernelType, class NormalizedKernelType, int Dims>
-  KernelType *ResetHostKernelHelper(const KernelType &KernelFunc) {
-    NormalizedKernelType NormalizedKernel(KernelFunc);
-    auto NormalizedKernelFunc =
-        std::function<void(const sycl::nd_item<Dims> &)>(NormalizedKernel);
-    auto HostKernelPtr = new detail::HostKernel<decltype(NormalizedKernelFunc),
-                                                sycl::nd_item<Dims>, Dims>(
-        std::move(NormalizedKernelFunc));
-    MHostKernel.reset(HostKernelPtr);
-    return &HostKernelPtr->MKernel.template target<NormalizedKernelType>()
-                ->MKernelFunc;
-  }
-
-  // For 'sycl::id<Dims>' kernel argument
-  template <class KernelType, typename ArgT, int Dims>
-  std::enable_if_t<std::is_same_v<ArgT, sycl::id<Dims>>, KernelType *>
-  ResetHostKernel(const KernelType &KernelFunc) {
-    struct NormalizedKernelType {
-      KernelType MKernelFunc;
-      NormalizedKernelType(const KernelType &KernelFunc)
-          : MKernelFunc(KernelFunc) {}
-      void operator()(const nd_item<Dims> &Arg) {
-        detail::runKernelWithArg(MKernelFunc, Arg.get_global_id());
-      }
-    };
-    return ResetHostKernelHelper<KernelType, struct NormalizedKernelType, Dims>(
-        KernelFunc);
-  }
-
-  // For 'sycl::nd_item<Dims>' kernel argument
-  template <class KernelType, typename ArgT, int Dims>
-  std::enable_if_t<std::is_same_v<ArgT, sycl::nd_item<Dims>>, KernelType *>
-  ResetHostKernel(const KernelType &KernelFunc) {
-    struct NormalizedKernelType {
-      KernelType MKernelFunc;
-      NormalizedKernelType(const KernelType &KernelFunc)
-          : MKernelFunc(KernelFunc) {}
-      void operator()(const nd_item<Dims> &Arg) {
-        detail::runKernelWithArg(MKernelFunc, Arg);
-      }
-    };
-    return ResetHostKernelHelper<KernelType, struct NormalizedKernelType, Dims>(
-        KernelFunc);
-  }
-
-  // For 'sycl::item<Dims, without_offset>' kernel argument
-  template <class KernelType, typename ArgT, int Dims>
-  std::enable_if_t<std::is_same_v<ArgT, sycl::item<Dims, false>>, KernelType *>
-  ResetHostKernel(const KernelType &KernelFunc) {
-    struct NormalizedKernelType {
-      KernelType MKernelFunc;
-      NormalizedKernelType(const KernelType &KernelFunc)
-          : MKernelFunc(KernelFunc) {}
-      void operator()(const nd_item<Dims> &Arg) {
-        sycl::item<Dims, false> Item = detail::Builder::createItem<Dims, false>(
-            Arg.get_global_range(), Arg.get_global_id());
-        detail::runKernelWithArg(MKernelFunc, Item);
-      }
-    };
-    return ResetHostKernelHelper<KernelType, struct NormalizedKernelType, Dims>(
-        KernelFunc);
-  }
-
-  // For 'sycl::item<Dims, with_offset>' kernel argument
-  template <class KernelType, typename ArgT, int Dims>
-  std::enable_if_t<std::is_same_v<ArgT, sycl::item<Dims, true>>, KernelType *>
-  ResetHostKernel(const KernelType &KernelFunc) {
-    struct NormalizedKernelType {
-      KernelType MKernelFunc;
-      NormalizedKernelType(const KernelType &KernelFunc)
-          : MKernelFunc(KernelFunc) {}
-      void operator()(const nd_item<Dims> &Arg) {
-        sycl::item<Dims, true> Item = detail::Builder::createItem<Dims, true>(
-            Arg.get_global_range(), Arg.get_global_id(), Arg.get_offset());
-        detail::runKernelWithArg(MKernelFunc, Item);
-      }
-    };
-    return ResetHostKernelHelper<KernelType, struct NormalizedKernelType, Dims>(
-        KernelFunc);
-  }
-
-  // For 'void' kernel argument (single_task)
-  template <class KernelType, typename ArgT, int Dims>
-  typename std::enable_if_t<std::is_same_v<ArgT, void>, KernelType *>
-  ResetHostKernel(const KernelType &KernelFunc) {
-    struct NormalizedKernelType {
-      KernelType MKernelFunc;
-      NormalizedKernelType(const KernelType &KernelFunc)
-          : MKernelFunc(KernelFunc) {}
-      void operator()(const nd_item<Dims> &Arg) {
-        (void)Arg;
-        detail::runKernelWithoutArg(MKernelFunc);
-      }
-    };
-    return ResetHostKernelHelper<KernelType, struct NormalizedKernelType, Dims>(
-        KernelFunc);
-  }
-
-  // For 'sycl::group<Dims>' kernel argument
-  // 'wrapper'-based approach using 'NormalizedKernelType' struct is not used
-  // for 'void(sycl::group<Dims>)' since 'void(sycl::group<Dims>)' is not
-  // supported in ESIMD.
-  template <class KernelType, typename ArgT, int Dims>
-  std::enable_if_t<std::is_same_v<ArgT, sycl::group<Dims>>, KernelType *>
-  ResetHostKernel(const KernelType &KernelFunc) {
-    MHostKernel.reset(
-        new detail::HostKernel<KernelType, ArgT, Dims>(KernelFunc));
-    return (KernelType *)(MHostKernel->getPtr());
-  }
-
   /// Verifies the kernel bundle to be used if any is set. This throws a
   /// sycl::exception with error code errc::kernel_not_supported if the used
   /// kernel bundle does not contain a suitable device image with the requested
@@ -918,8 +794,8 @@ class __SYCL_EXPORT handler {
         detail::KernelLambdaHasKernelHandlerArgT<KernelType,
                                                  LambdaArgType>::value;
 
-    KernelType *KernelPtr =
-        ResetHostKernel<KernelType, LambdaArgType, Dims>(KernelFunc);
+    MHostKernel.reset(
+        new detail::HostKernel<KernelType, LambdaArgType, Dims>(KernelFunc));
 
     constexpr bool KernelHasName =
         detail::getKernelName<KernelName>() != nullptr &&
@@ -950,7 +826,7 @@ class __SYCL_EXPORT handler {
     if (KernelHasName) {
       // TODO support ESIMD in no-integration-header case too.
       clearArgs();
-      extractArgsAndReqsFromLambda(reinterpret_cast<char *>(KernelPtr),
+      extractArgsAndReqsFromLambda(MHostKernel->getPtr(),
                                    detail::getKernelParamDescs<KernelName>(),
                                    detail::isKernelESIMD<KernelName>());
       MKernelName = detail::getKernelName<KernelName>();