[SYCL][ESIMD][EMU] ESIMD_CPU Kernel launch and ESIMD_EMU backend loading (#4020)

* [SYCL][ESIMD][EMU] ESIMD_CPU Kernel launch and Emulated Intrinsics

* This PR is for enabling kernel launching for ESIMD_CPU
* esimd_cpu backend is loaded in SYCL runtime
* Base PR : #4011

Author: dongkyunahn-intel <[email protected]>

Author: dongkyunahn-intel

sycl/CMakeLists.txt

@@ -326,17 +326,17 @@ if(SYCL_BUILD_PI_HIP)
   list(APPEND SYCL_TOOLCHAIN_DEPLOY_COMPONENTS libspirv-builtins pi_hip)
 endif()
 
-# TODO : Remove 'if (NOT MSVC)' when CM_EMU supports Windows
-# environment
-if (NOT MSVC)
-  if (SYCL_BUILD_PI_ESIMD_EMULATOR)
-    list(APPEND SYCL_TOOLCHAIN_DEPLOY_COMPONENTS pi_esimd_emulator libcmrt-headers)
-    if (MSVC)
-      list(APPEND SYCL_TOOLCHAIN_DEPLOY_COMPONENTS libcmrt-libs libcmrt-dlls)
-    else()
-      list(APPEND SYCL_TOOLCHAIN_DEPLOY_COMPONENTS libcmrt-sos)
-    endif()
+if (SYCL_BUILD_PI_ESIMD_EMULATOR)
+  list(APPEND SYCL_TOOLCHAIN_DEPLOY_COMPONENTS pi_esimd_emulator libcmrt-headers)
+  if (MSVC)
+    list(APPEND SYCL_TOOLCHAIN_DEPLOY_COMPONENTS libcmrt-libs libcmrt-dlls)
+  else()
+    list(APPEND SYCL_TOOLCHAIN_DEPLOY_COMPONENTS libcmrt-sos)
   endif()
+else()
+  # TODO/FIXME : Removing empty header file (cm_rt.h) generation when
+  # the ESIMD_EMULATOR support is enabled by default
+  file (TOUCH ${SYCL_INCLUDE_BUILD_DIR}/sycl/CL/cm_rt.h)
 endif()
 
 # Use it as fake dependency in order to force another command(s) to execute.

sycl/include/CL/sycl/detail/cg_types.hpp

@@ -246,6 +246,9 @@ template <class KernelType, class KernelArgType, int Dims, typename KernelName>
 class HostKernel : public HostKernelBase {
   using IDBuilder = sycl::detail::Builder;
   KernelType MKernel;
+  // Allowing accessing MKernel from 'ResetHostKernelHelper' method of
+  // 'sycl::handler'
+  friend class sycl::handler;
 
 public:
   HostKernel(KernelType Kernel) : MKernel(Kernel) {}

sycl/include/CL/sycl/handler.hpp

@@ -561,6 +561,124 @@ class __SYCL_EXPORT handler {
     }
   }
 
+  /* The kernel passed to StoreLambda can take an id, an item or an nd_item as
+   * its argument. Since esimd plugin directly invokes the kernel (doesn’t use
+   * piKernelSetArg), the kernel argument type must be known to the plugin.
+   * However, passing kernel argument type to the plugin 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 plugin.
+   */
+
+  // For 'id, item w/wo offset, nd_item' kernel arguments
+  template <class KernelType, class NormalizedKernelType, int Dims,
+            typename KernelName>
+  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, KernelName>(
+            NormalizedKernelFunc);
+    MHostKernel.reset(HostKernelPtr);
+    return &HostKernelPtr->MKernel.template target<NormalizedKernelType>()
+                ->MKernelFunc;
+  }
+
+  // For 'sycl::id<Dims>' kernel argument
+  template <class KernelType, typename ArgT, int Dims, typename KernelName>
+  typename std::enable_if<std::is_same<ArgT, sycl::id<Dims>>::value,
+                          KernelType *>::type
+  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,
+                                 KernelName>(KernelFunc);
+  }
+
+  // For 'sycl::nd_item<Dims>' kernel argument
+  template <class KernelType, typename ArgT, int Dims, typename KernelName>
+  typename std::enable_if<std::is_same<ArgT, sycl::nd_item<Dims>>::value,
+                          KernelType *>::type
+  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,
+                                 KernelName>(KernelFunc);
+  }
+
+  // For 'sycl::item<Dims, without_offset>' kernel argument
+  template <class KernelType, typename ArgT, int Dims, typename KernelName>
+  typename std::enable_if<std::is_same<ArgT, sycl::item<Dims, false>>::value,
+                          KernelType *>::type
+  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,
+                                 KernelName>(KernelFunc);
+  }
+
+  // For 'sycl::item<Dims, with_offset>' kernel argument
+  template <class KernelType, typename ArgT, int Dims, typename KernelName>
+  typename std::enable_if<std::is_same<ArgT, sycl::item<Dims, true>>::value,
+                          KernelType *>::type
+  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,
+                                 KernelName>(KernelFunc);
+  }
+
+  /* 'wrapper'-based approach using 'NormalizedKernelType' struct is
+   * not applied for 'void(void)' type kernel and
+   * 'void(sycl::group<Dims>)'. This is because 'void(void)' type does
+   * not have argument to normalize and 'void(sycl::group<Dims>)' is
+   * not supported in ESIMD.
+   */
+  // For 'void' and 'sycl::group<Dims>' kernel argument
+  template <class KernelType, typename ArgT, int Dims, typename KernelName>
+  typename std::enable_if<std::is_same<ArgT, void>::value ||
+                              std::is_same<ArgT, sycl::group<Dims>>::value,
+                          KernelType *>::type
+  ResetHostKernel(const KernelType &KernelFunc) {
+    MHostKernel.reset(
+        new detail::HostKernel<KernelType, ArgT, Dims, KernelName>(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
@@ -588,18 +706,19 @@ class __SYCL_EXPORT handler {
           "kernel_handler is not yet supported by host device.",
           PI_INVALID_OPERATION);
     }
-    MHostKernel.reset(
-        new detail::HostKernel<KernelType, LambdaArgType, Dims, KernelName>(
-            KernelFunc));
+    KernelType *KernelPtr =
+        ResetHostKernel<KernelType, LambdaArgType, Dims, KernelName>(
+            KernelFunc);
 
     using KI = sycl::detail::KernelInfo<KernelName>;
     // Empty name indicates that the compilation happens without integration
     // header, so don't perform things that require it.
     if (KI::getName() != nullptr && KI::getName()[0] != '\0') {
       // TODO support ESIMD in no-integration-header case too.
       MArgs.clear();
-      extractArgsAndReqsFromLambda(MHostKernel->getPtr(), KI::getNumParams(),
-                                   &KI::getParamDesc(0), KI::isESIMD());
+      extractArgsAndReqsFromLambda(reinterpret_cast<char *>(KernelPtr),
+                                   KI::getNumParams(), &KI::getParamDesc(0),
+                                   KI::isESIMD());
       MKernelName = KI::getName();
       MOSModuleHandle = detail::OSUtil::getOSModuleHandle(KI::getName());
     } else {

sycl/include/sycl/ext/intel/experimental/esimd/detail/atomic_intrin.hpp

@@ -0,0 +1,21 @@
+//==-------- atomic_intrin.hpp - Atomic intrinsic definition file ----------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+#pragma once
+
+#include <CL/sycl/exception.hpp>
+
+// This function implements atomic update of pre-existing variable in the
+// absense of C++ 20's atomic_ref.
+template <typename Ty> Ty atomic_add_fetch(Ty *ptr, Ty val) {
+#ifdef _WIN32
+  // TODO: Windows will be supported soon
+  throw cl::sycl::feature_not_supported();
+#else
+  return __atomic_add_fetch(ptr, val, __ATOMIC_RELAXED);
+#endif
+}

sycl/include/sycl/ext/intel/experimental/esimd/emu/detail/esimd_emulator_device_interface.hpp

@@ -65,56 +65,8 @@ struct ESIMDEmuPluginOpaqueData {
   uintptr_t version;
   void *data;
 };
-// The table below shows the correspondence between the \c version
-// and the contents of the \c data field:
-// version == 0, data is ESIMDDeviceInterface*
 
-ESIMDDeviceInterface *getESIMDDeviceInterface() {
-  // TODO (performance) cache the interface pointer, can make a difference
-  // when calling fine-grained libCM APIs through it (like memory access in a
-  // tight loop)
-  void *PIOpaqueData = nullptr;
-
-  PIOpaqueData =
-      getPluginOpaqueData<cl::sycl::backend::ext_intel_esimd_emulator>(nullptr);
-
-  ESIMDEmuPluginOpaqueData *OpaqueData =
-      reinterpret_cast<ESIMDEmuPluginOpaqueData *>(PIOpaqueData);
-
-  // First check if opaque data version is compatible.
-  if (OpaqueData->version != ESIMD_EMULATOR_PLUGIN_OPAQUE_DATA_VERSION) {
-    // NOTE: the version check should always be '!=' as layouts of different
-    // versions of PluginOpaqueData is not backward compatible, unlike
-    // layout of the ESIMDDeviceInterface.
-
-    std::cerr << __FUNCTION__ << std::endl
-              << "Opaque data returned by ESIMD Emu plugin is incompatible with"
-              << "the one used in current implementation." << std::endl
-              << "Returned version : " << OpaqueData->version << std::endl
-              << "Required version : "
-              << ESIMD_EMULATOR_PLUGIN_OPAQUE_DATA_VERSION << std::endl;
-    throw cl::sycl::feature_not_supported();
-  }
-  // Opaque data version is OK, can cast the 'data' field.
-  ESIMDDeviceInterface *Interface =
-      reinterpret_cast<ESIMDDeviceInterface *>(OpaqueData->data);
-
-  // Now check that device interface version is compatible.
-  if (Interface->version < ESIMD_DEVICE_INTERFACE_VERSION) {
-    std::cerr << __FUNCTION__ << std::endl
-              << "The device interface version provided from plug-in "
-              << "library is behind required device interface version"
-              << std::endl
-              << "Found version : " << Interface->version << std::endl
-              << "Required version :" << ESIMD_DEVICE_INTERFACE_VERSION
-              << std::endl;
-    throw cl::sycl::feature_not_supported();
-  }
-  return Interface;
-}
-
-#undef ESIMD_DEVICE_INTERFACE_VERSION
-#undef ESIMD_EMULATOR_PLUGIN_OPAQUE_DATA_VERSION
+__SYCL_EXPORT ESIMDDeviceInterface *getESIMDDeviceInterface();
 
 } // namespace detail
 } // namespace sycl

sycl/source/CMakeLists.txt

@@ -176,6 +176,7 @@ set(SYCL_SOURCES
     "sampler.cpp"
     "stream.cpp"
     "spirv_ops.cpp"
+    "esimd_emulator_device_interface.cpp"
     "$<$<PLATFORM_ID:Windows>:detail/windows_pi.cpp>"
     "$<$<OR:$<PLATFORM_ID:Linux>,$<PLATFORM_ID:Darwin>>:detail/posix_pi.cpp>"
     "$<$<PLATFORM_ID:Windows>:abi_replacements_windows.cpp>"

sycl/source/detail/config.cpp

@@ -169,13 +169,14 @@ getSyclDeviceTypeMap() {
 }
 
 // Array is used by SYCL_DEVICE_FILTER and SYCL_DEVICE_ALLOWLIST
-const std::array<std::pair<std::string, backend>, 6> &getSyclBeMap() {
-  static const std::array<std::pair<std::string, backend>, 6> SyclBeMap = {
+const std::array<std::pair<std::string, backend>, 7> &getSyclBeMap() {
+  static const std::array<std::pair<std::string, backend>, 7> SyclBeMap = {
       {{"host", backend::host},
        {"opencl", backend::opencl},
        {"level_zero", backend::ext_oneapi_level_zero},
        {"cuda", backend::ext_oneapi_cuda},
        {"hip", backend::ext_oneapi_hip},
+       {"esimd_emulator", backend::ext_intel_esimd_emulator},
        {"*", backend::all}}};
   return SyclBeMap;
 }

sycl/source/detail/config.hpp

@@ -134,12 +134,13 @@ template <> class SYCLConfig<SYCL_BE> {
       return BackendPtr;
 
     const char *ValStr = BaseT::getRawValue();
-    const std::array<std::pair<std::string, backend>, 5> SyclBeMap = {
+    const std::array<std::pair<std::string, backend>, 6> SyclBeMap = {
         {{"PI_OPENCL", backend::opencl},
          {"PI_LEVEL_ZERO", backend::ext_oneapi_level_zero},
          {"PI_LEVEL0", backend::ext_oneapi_level_zero}, // for backward
                                                         // compatibility
          {"PI_CUDA", backend::ext_oneapi_cuda},
+         {"PI_ESIMD_EMULATOR", backend::ext_intel_esimd_emulator},
          {"PI_HIP", backend::ext_oneapi_hip}}};
     if (ValStr) {
       auto It = std::find_if(
@@ -149,7 +150,8 @@ template <> class SYCLConfig<SYCL_BE> {
           });
       if (It == SyclBeMap.end())
         pi::die("Invalid backend. "
-                "Valid values are PI_OPENCL/PI_LEVEL_ZERO/PI_CUDA/PI_HIP");
+                "Valid values are "
+                "PI_OPENCL/PI_LEVEL_ZERO/PI_CUDA/PI_ESIMD_EMULATOR/PI_HIP");
       static backend Backend = It->second;
       BackendPtr = &Backend;
     }
@@ -243,7 +245,7 @@ const std::array<std::pair<std::string, info::device_type>, 5> &
 getSyclDeviceTypeMap();
 
 // Array is used by SYCL_DEVICE_FILTER and SYCL_DEVICE_ALLOWLIST
-const std::array<std::pair<std::string, backend>, 6> &getSyclBeMap();
+const std::array<std::pair<std::string, backend>, 7> &getSyclBeMap();
 
 template <> class SYCLConfig<SYCL_DEVICE_FILTER> {
   using BaseT = SYCLConfigBase<SYCL_DEVICE_FILTER>;

sycl/source/detail/device_filter.cpp

@@ -91,7 +91,7 @@ device_filter::device_filter(const std::string &FilterString) {
       std::string Message =
           std::string("Invalid device filter: ") + FilterString +
           "\nPossible backend values are "
-          "{host,opencl,level_zero,cuda,hip,*}.\n"
+          "{host,opencl,level_zero,cuda,hip,esimd_emulator*}.\n"
           "Possible device types are {host,cpu,gpu,acc,*}.\n"
           "Device number should be an non-negative integer.\n";
       throw cl::sycl::invalid_parameter_error(Message, PI_INVALID_VALUE);

sycl/source/detail/pi.cpp

@@ -285,11 +285,14 @@ std::vector<std::pair<std::string, backend>> findPlugins() {
                              backend::ext_oneapi_level_zero);
     PluginNames.emplace_back(__SYCL_CUDA_PLUGIN_NAME, backend::ext_oneapi_cuda);
     PluginNames.emplace_back(__SYCL_HIP_PLUGIN_NAME, backend::ext_oneapi_hip);
+    PluginNames.emplace_back(__SYCL_ESIMD_EMULATOR_PLUGIN_NAME,
+                             backend::ext_intel_esimd_emulator);
   } else {
     std::vector<device_filter> Filters = FilterList->get();
     bool OpenCLFound = false;
     bool LevelZeroFound = false;
     bool CudaFound = false;
+    bool EsimdCpuFound = false;
     bool HIPFound = false;
     for (const device_filter &Filter : Filters) {
       backend Backend = Filter.Backend;
@@ -310,6 +313,12 @@ std::vector<std::pair<std::string, backend>> findPlugins() {
                                  backend::ext_oneapi_cuda);
         CudaFound = true;
       }
+      if (!EsimdCpuFound && (Backend == backend::ext_intel_esimd_emulator ||
+                             Backend == backend::all)) {
+        PluginNames.emplace_back(__SYCL_ESIMD_EMULATOR_PLUGIN_NAME,
+                                 backend::ext_intel_esimd_emulator);
+        EsimdCpuFound = true;
+      }
       if (!HIPFound &&
           (Backend == backend::ext_oneapi_hip || Backend == backend::all)) {
         PluginNames.emplace_back(__SYCL_HIP_PLUGIN_NAME,
@@ -429,6 +438,12 @@ static void initializePlugins(std::vector<plugin> &Plugins) {
       // Use the LEVEL_ZERO plugin as the GlobalPlugin
       GlobalPlugin = std::make_shared<plugin>(
           PluginInformation, backend::ext_oneapi_level_zero, Library);
+    } else if (InteropBE == backend::ext_intel_esimd_emulator &&
+               PluginNames[I].first.find("esimd_emulator") !=
+                   std::string::npos) {
+      // Use the ESIMD_EMULATOR plugin as the GlobalPlugin
+      GlobalPlugin = std::make_shared<plugin>(
+          PluginInformation, backend::ext_intel_esimd_emulator, Library);
     }
     Plugins.emplace_back(
         plugin(PluginInformation, PluginNames[I].second, Library));