[SYCL] Remove ESIMD Emulator

sycl/doc/GetStartedGuide.md

@@ -12,7 +12,6 @@ and a wide range of compute accelerators such as GPU and FPGA.
     * [Build DPC++ toolchain with support for NVIDIA CUDA](#build-dpc-toolchain-with-support-for-nvidia-cuda)
     * [Build DPC++ toolchain with support for HIP AMD](#build-dpc-toolchain-with-support-for-hip-amd)
     * [Build DPC++ toolchain with support for HIP NVIDIA](#build-dpc-toolchain-with-support-for-hip-nvidia)
-    * [Build DPC++ toolchain with support for ESIMD CPU Emulation](#build-dpc-toolchain-with-support-for-esimd-cpu-emulation)
     * [Build DPC++ toolchain with support for runtime kernel fusion](#build-dpc-toolchain-with-support-for-runtime-kernel-fusion)
     * [Build DPC++ toolchain with a custom Unified Runtime](#build-dpc-toolchain-with-a-custom-unified-runtime)
     * [Build Doxygen documentation](#build-doxygen-documentation)
@@ -302,32 +301,6 @@ as well as the CUDA Runtime API to be installed, see
 Currently, this has only been tried on Linux, with ROCm 4.2.0 or 4.3.0, with
 CUDA 11, and using a GeForce 1060 device.
 
-### Build DPC++ toolchain with support for ESIMD CPU Emulation
-
-There is experimental support for DPC++ for using ESIMD CPU Emulation.
-
-This feature supports ESIMD CPU Emulation using CM_EMU library
-[CM Emulation project](https://github.com/intel/cm-cpu-emulation). The library
-package will be generated from source codes downloaded from its open source
-project and installed in your deploy directory during toolchain build.
-
-To enable support for ESIMD CPU emulation, follow the instructions for the Linux
-DPC++ toolchain, but add the `--enable-esimd-emulator`.
-
-Enabling this flag requires following packages installed.
-
-* Ubuntu 22.04
-  * libva-dev / 2.7.0-2
-  * libffi-dev / 3.3-4
-  * libtool
-* RHEL 8.\*
-  * libffi
-  * libffi-devel
-  * libva
-  * libva-devel
-
-Currently, this feature was tested and verified on Ubuntu 22.04 environment.
-
 ### Build DPC++ toolchain with support for runtime kernel fusion
 
 Support for the experimental SYCL extension for user-driven kernel fusion at

sycl/include/sycl/backend_types.hpp

@@ -21,8 +21,8 @@ enum class backend : char {
   ext_oneapi_level_zero = 2,
   ext_oneapi_cuda = 3,
   all = 4,
-  ext_intel_esimd_emulator __SYCL_DEPRECATED(
-      "esimd emulator is no longer supported") = 5,
+  // No support anymore:
+  // ext_intel_esimd_emulator  = 5,
   ext_oneapi_hip = 6,
   ext_oneapi_native_cpu = 7,
 };
@@ -50,9 +50,6 @@ inline std::ostream &operator<<(std::ostream &Out, backend be) {
   case backend::ext_oneapi_cuda:
     Out << "ext_oneapi_cuda";
     break;
-  case backend::ext_intel_esimd_emulator:
-    Out << "ext_intel_esimd_emulator";
-    break;
   case backend::ext_oneapi_hip:
     Out << "ext_oneapi_hip";
     break;
@@ -76,8 +73,6 @@ inline std::string_view get_backend_name_no_vendor(backend Backend) {
     return "level_zero";
   case backend::ext_oneapi_cuda:
     return "cuda";
-  case backend::ext_intel_esimd_emulator:
-    return "esimd_emulator";
   case backend::ext_oneapi_hip:
     return "hip";
   case backend::ext_oneapi_native_cpu:

sycl/include/sycl/detail/pi.h

@@ -312,7 +312,8 @@ typedef enum {
   PI_EXT_PLATFORM_BACKEND_OPENCL = 2,     ///< The backend is OpenCL
   PI_EXT_PLATFORM_BACKEND_CUDA = 3,       ///< The backend is CUDA
   PI_EXT_PLATFORM_BACKEND_HIP = 4,        ///< The backend is HIP
-  PI_EXT_PLATFORM_BACKEND_ESIMD = 5,      ///< The backend is ESIMD
+  // Not supported anymore:
+  // PI_EXT_PLATFORM_BACKEND_ESIMD = 5,
   PI_EXT_PLATFORM_BACKEND_NATIVE_CPU = 6, ///< The backend is NATIVE_CPU
 } _pi_platform_backend;
 

sycl/include/sycl/detail/pi.hpp

@@ -69,15 +69,13 @@ bool trace(TraceLevel level);
 #define __SYCL_OPENCL_PLUGIN_NAME "pi_opencl.dll"
 #define __SYCL_LEVEL_ZERO_PLUGIN_NAME "pi_level_zero.dll"
 #define __SYCL_CUDA_PLUGIN_NAME "pi_cuda.dll"
-#define __SYCL_ESIMD_EMULATOR_PLUGIN_NAME "pi_esimd_emulator.dll"
 #define __SYCL_HIP_PLUGIN_NAME "pi_hip.dll"
 #define __SYCL_UR_PLUGIN_NAME "pi_unified_runtime.dll"
 #define __SYCL_NATIVE_CPU_PLUGIN_NAME "pi_native_cpu.dll"
 #else
 #define __SYCL_OPENCL_PLUGIN_NAME "libpi_opencl.dll"
 #define __SYCL_LEVEL_ZERO_PLUGIN_NAME "libpi_level_zero.dll"
 #define __SYCL_CUDA_PLUGIN_NAME "libpi_cuda.dll"
-#define __SYCL_ESIMD_EMULATOR_PLUGIN_NAME "libpi_esimd_emulator.dll"
 #define __SYCL_HIP_PLUGIN_NAME "libpi_hip.dll"
 #define __SYCL_UR_PLUGIN_NAME "libpi_unified_runtime.dll"
 #define __SYCL_NATIVE_CPU_PLUGIN_NAME "libpi_native_cpu.dll"
@@ -86,15 +84,13 @@ bool trace(TraceLevel level);
 #define __SYCL_OPENCL_PLUGIN_NAME "libpi_opencl.so"
 #define __SYCL_LEVEL_ZERO_PLUGIN_NAME "libpi_level_zero.so"
 #define __SYCL_CUDA_PLUGIN_NAME "libpi_cuda.so"
-#define __SYCL_ESIMD_EMULATOR_PLUGIN_NAME "libpi_esimd_emulator.so"
 #define __SYCL_HIP_PLUGIN_NAME "libpi_hip.so"
 #define __SYCL_UR_PLUGIN_NAME "libpi_unified_runtime.so"
 #define __SYCL_NATIVE_CPU_PLUGIN_NAME "libpi_native_cpu.so"
 #elif defined(__SYCL_RT_OS_DARWIN)
 #define __SYCL_OPENCL_PLUGIN_NAME "libpi_opencl.dylib"
 #define __SYCL_LEVEL_ZERO_PLUGIN_NAME "libpi_level_zero.dylib"
 #define __SYCL_CUDA_PLUGIN_NAME "libpi_cuda.dylib"
-#define __SYCL_ESIMD_EMULATOR_PLUGIN_NAME "libpi_esimd_emulator.dylib"
 #define __SYCL_HIP_PLUGIN_NAME "libpi_hip.dylib"
 #define __SYCL_UR_PLUGIN_NAME "libpi_unified_runtime.dylib"
 #define __SYCL_NATIVE_CPU_PLUGIN_NAME "libpi_native_cpu.dylib"

sycl/pi_win_proxy_loader/pi_win_proxy_loader.cpp

@@ -89,15 +89,13 @@ std::wstring getCurrentDSODir() {
 #define __SYCL_OPENCL_PLUGIN_NAME "pi_opencl.dll"
 #define __SYCL_LEVEL_ZERO_PLUGIN_NAME "pi_level_zero.dll"
 #define __SYCL_CUDA_PLUGIN_NAME "pi_cuda.dll"
-#define __SYCL_ESIMD_EMULATOR_PLUGIN_NAME "pi_esimd_emulator.dll"
 #define __SYCL_HIP_PLUGIN_NAME "pi_hip.dll"
 #define __SYCL_UNIFIED_RUNTIME_PLUGIN_NAME "pi_unified_runtime.dll"
 #define __SYCL_NATIVE_CPU_PLUGIN_NAME "pi_native_cpu.dll"
 #else // llvm-mingw
 #define __SYCL_OPENCL_PLUGIN_NAME "libpi_opencl.dll"
 #define __SYCL_LEVEL_ZERO_PLUGIN_NAME "libpi_level_zero.dll"
 #define __SYCL_CUDA_PLUGIN_NAME "libpi_cuda.dll"
-#define __SYCL_ESIMD_EMULATOR_PLUGIN_NAME "libpi_esimd_emulator.dll"
 #define __SYCL_HIP_PLUGIN_NAME "libpi_hip.dll"
 #define __SYCL_UNIFIED_RUNTIME_PLUGIN_NAME "libpi_unified_runtime.dll"
 #define __SYCL_NATIVE_CPU_PLUGIN_NAME "libpi_native_cpu.dll"
@@ -147,7 +145,6 @@ void preloadLibraries() {
   loadPlugin(__SYCL_OPENCL_PLUGIN_NAME);
   loadPlugin(__SYCL_LEVEL_ZERO_PLUGIN_NAME);
   loadPlugin(__SYCL_CUDA_PLUGIN_NAME);
-  loadPlugin(__SYCL_ESIMD_EMULATOR_PLUGIN_NAME);
   loadPlugin(__SYCL_HIP_PLUGIN_NAME);
   loadPlugin(__SYCL_UNIFIED_RUNTIME_PLUGIN_NAME);
   loadPlugin(__SYCL_NATIVE_CPU_PLUGIN_NAME);

sycl/source/CMakeLists.txt

@@ -261,7 +261,6 @@ set(SYCL_COMMON_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>"
 )

sycl/source/backend.cpp

@@ -58,8 +58,6 @@ backend convertBackend(pi_platform_backend PiBackend) {
     return backend::ext_oneapi_cuda;
   case PI_EXT_PLATFORM_BACKEND_HIP:
     return backend::ext_oneapi_hip;
-  case PI_EXT_PLATFORM_BACKEND_ESIMD:
-    return backend::ext_intel_esimd_emulator;
   case PI_EXT_PLATFORM_BACKEND_NATIVE_CPU:
     return backend::ext_oneapi_native_cpu;
   }

sycl/source/detail/config.cpp

@@ -163,17 +163,15 @@ void dumpConfig() {
 
 // Array is used by SYCL_DEVICE_FILTER and SYCL_DEVICE_ALLOWLIST and
 // ONEAPI_DEVICE_SELECTOR
-// TODO: Remove esimd_emulator in the next ABI breaking window.
 // TODO: host device type will be removed once sycl_ext_oneapi_filter_selector
 // is removed.
-const std::array<std::pair<std::string, backend>, 8> &getSyclBeMap() {
-  static const std::array<std::pair<std::string, backend>, 8> 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},
        {"native_cpu", backend::ext_oneapi_native_cpu},
        {"*", backend::all}}};
   return SyclBeMap;

sycl/source/detail/config.hpp

@@ -251,7 +251,7 @@ getSyclDeviceTypeMap() {
 
 // Array is used by SYCL_DEVICE_FILTER and SYCL_DEVICE_ALLOWLIST and
 // ONEAPI_DEVICE_SELECTOR
-const std::array<std::pair<std::string, backend>, 8> &getSyclBeMap();
+const std::array<std::pair<std::string, backend>, 7> &getSyclBeMap();
 
 // ---------------------------------------
 // ONEAPI_DEVICE_SELECTOR support

sycl/source/detail/device_impl.cpp

@@ -363,7 +363,7 @@ bool device_impl::has(aspect Aspect) const {
     return is_accelerator();
   case aspect::custom:
     return false;
-  // TODO: Implement this for FPGA and ESIMD emulators.
+  // TODO: Implement this for FPGA emulator.
   case aspect::emulated:
     return false;
   case aspect::host_debuggable:

sycl/source/detail/pi.cpp

@@ -68,9 +68,6 @@ template <sycl::backend BE> void *getPluginOpaqueData(void *OpaqueDataParam) {
   return ReturnOpaqueData;
 }
 
-template __SYCL_EXPORT void *
-getPluginOpaqueData<sycl::backend::ext_intel_esimd_emulator>(void *);
-
 namespace pi {
 
 static void initializePlugins(std::vector<PluginPtr> &Plugins);
@@ -514,8 +511,6 @@ template <backend BE> const PluginPtr &getPlugin() {
 template __SYCL_EXPORT const PluginPtr &getPlugin<backend::opencl>();
 template __SYCL_EXPORT const PluginPtr &
 getPlugin<backend::ext_oneapi_level_zero>();
-template __SYCL_EXPORT const PluginPtr &
-getPlugin<backend::ext_intel_esimd_emulator>();
 template __SYCL_EXPORT const PluginPtr &getPlugin<backend::ext_oneapi_cuda>();
 template __SYCL_EXPORT const PluginPtr &getPlugin<backend::ext_oneapi_hip>();
 

sycl/source/detail/scheduler/commands.cpp

@@ -2992,8 +2992,7 @@ pi_int32 ExecCGCommand::enqueueImpQueue() {
     NDRDescT &NDRDesc = ExecKernel->MNDRDesc;
     std::vector<ArgDesc> &Args = ExecKernel->MArgs;
 
-    if (MQueue->is_host() || (MQueue->getDeviceImplPtr()->getBackend() ==
-                              backend::ext_intel_esimd_emulator)) {
+    if (MQueue->is_host()) {
       for (ArgDesc &Arg : Args)
         if (kernel_param_kind_t::kind_accessor == Arg.MType) {
           Requirement *Req = (Requirement *)(Arg.MPtr);
@@ -3006,20 +3005,8 @@ pi_int32 ExecCGCommand::enqueueImpQueue() {
         Plugin->call<PiApiKind::piEventsWait>(RawEvents.size(), &RawEvents[0]);
       }
 
-      if (MQueue->is_host()) {
-        ExecKernel->MHostKernel->call(NDRDesc,
-                                      getEvent()->getHostProfilingInfo());
-      } else {
-        assert(MQueue->getDeviceImplPtr()->getBackend() ==
-               backend::ext_intel_esimd_emulator);
-        if (MEvent != nullptr)
-          MEvent->setHostEnqueueTime();
-        MQueue->getPlugin()->call<PiApiKind::piEnqueueKernelLaunch>(
-            nullptr,
-            reinterpret_cast<pi_kernel>(ExecKernel->MHostKernel->getPtr()),
-            NDRDesc.Dims, &NDRDesc.GlobalOffset[0], &NDRDesc.GlobalSize[0],
-            &NDRDesc.LocalSize[0], 0, nullptr, nullptr);
-      }
+      ExecKernel->MHostKernel->call(NDRDesc,
+                                    getEvent()->getHostProfilingInfo());
       return PI_SUCCESS;
     }
 

sycl/source/device_selector.cpp

@@ -179,13 +179,6 @@ __SYCL_EXPORT int default_selector_v(const device &dev) {
   // The default selector doesn't reject any devices.
   int Score = 0;
 
-  // we give the esimd_emulator device a score of zero to prevent it from being
-  // chosen among other devices. The same thing is done for gpu_selector_v
-  // below.
-  if (dev.get_backend() == backend::ext_intel_esimd_emulator) {
-    return 0;
-  }
-
   traceDeviceSelector("info::device_type::automatic");
 
   if (dev.is_gpu())
@@ -209,10 +202,6 @@ __SYCL_EXPORT int default_selector_v(const device &dev) {
 __SYCL_EXPORT int gpu_selector_v(const device &dev) {
   int Score = detail::REJECT_DEVICE_SCORE;
 
-  if (dev.get_backend() == backend::ext_intel_esimd_emulator) {
-    return 0;
-  }
-
   traceDeviceSelector("info::device_type::gpu");
   if (dev.is_gpu()) {
     Score = 1000;

sycl/source/handler.cpp

@@ -279,40 +279,10 @@ event handler::finalize() {
                                           : nullptr);
           Result = PI_SUCCESS;
         } else {
-          if (MQueue->getDeviceImplPtr()->getBackend() ==
-              backend::ext_intel_esimd_emulator) {
-            // Capture the host timestamp for profiling (queue time)
-            if (NewEvent != nullptr)
-              NewEvent->setHostEnqueueTime();
-            [&](auto... Args) {
-              if (MImpl->MKernelIsCooperative) {
-                MQueue->getPlugin()
-                    ->call<
-                        detail::PiApiKind::piextEnqueueCooperativeKernelLaunch>(
-                        Args...);
-              } else {
-                MQueue->getPlugin()
-                    ->call<detail::PiApiKind::piEnqueueKernelLaunch>(Args...);
-              }
-            }(/* queue */
-              nullptr,
-              /* kernel */
-              reinterpret_cast<pi_kernel>(MHostKernel->getPtr()),
-              /* work_dim */
-              MNDRDesc.Dims,
-              /* global_work_offset */ &MNDRDesc.GlobalOffset[0],
-              /* global_work_size */ &MNDRDesc.GlobalSize[0],
-              /* local_work_size */ &MNDRDesc.LocalSize[0],
-              /* num_events_in_wait_list */ 0,
-              /* event_wait_list */ nullptr,
-              /* event */ nullptr);
-            Result = PI_SUCCESS;
-          } else {
-            Result = enqueueImpKernel(
-                MQueue, MNDRDesc, MArgs, KernelBundleImpPtr, MKernel,
-                MKernelName.c_str(), RawEvents, NewEvent, nullptr,
-                MImpl->MKernelCacheConfig, MImpl->MKernelIsCooperative);
-          }
+          Result = enqueueImpKernel(
+              MQueue, MNDRDesc, MArgs, KernelBundleImpPtr, MKernel,
+              MKernelName.c_str(), RawEvents, NewEvent, nullptr,
+              MImpl->MKernelCacheConfig, MImpl->MKernelIsCooperative);
         }
 #ifdef XPTI_ENABLE_INSTRUMENTATION
         // Emit signal only when event is created

sycl/source/kernel_bundle.cpp

@@ -312,11 +312,6 @@ bool is_compatible(const std::vector<kernel_id> &KernelIDs, const device &Dev) {
                                        const detail::RTDeviceBinaryImage &Img) {
     const char *Target = Img.getRawData().DeviceTargetSpec;
     auto BE = Dev.get_backend();
-    // ESIMD emulator backend is only compatible with esimd kernels.
-    if (BE == sycl::backend::ext_intel_esimd_emulator) {
-      pi_device_binary_property Prop = Img.getProperty("isEsimdImage");
-      return (Prop && (detail::DeviceBinaryProperty(Prop).asUint32() != 0));
-    }
     if (strcmp(Target, __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64) == 0) {
       return (BE == sycl::backend::opencl ||
               BE == sycl::backend::ext_oneapi_level_zero);

sycl/test/abi/sycl_symbols_linux.dump

@@ -3359,7 +3359,6 @@ _ZN4sycl3_V16detail18stringifyErrorCodeEi
 _ZN4sycl3_V16detail19convertChannelOrderE23_pi_image_channel_order
 _ZN4sycl3_V16detail19convertChannelOrderENS0_19image_channel_orderE
 _ZN4sycl3_V16detail19getImageElementSizeEhNS0_18image_channel_typeE
-_ZN4sycl3_V16detail19getPluginOpaqueDataILNS0_7backendE5EEEPvS4_
 _ZN4sycl3_V16detail19kernel_bundle_plain21ext_oneapi_get_kernelERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
 _ZN4sycl3_V16detail19kernel_bundle_plain21ext_oneapi_has_kernelERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE
 _ZN4sycl3_V16detail19kernel_bundle_plain32set_specialization_constant_implEPKcPvm
@@ -3384,7 +3383,6 @@ _ZN4sycl3_V16detail22has_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6devi
 _ZN4sycl3_V16detail22reduGetPreferredWGSizeERSt10shared_ptrINS1_10queue_implEEm
 _ZN4sycl3_V16detail22removeDuplicateDevicesERKSt6vectorINS0_6deviceESaIS3_EE
 _ZN4sycl3_V16detail23constructorNotificationEPvS2_NS0_6access6targetENS3_4modeERKNS1_13code_locationE
-_ZN4sycl3_V16detail23getESIMDDeviceInterfaceEv
 _ZN4sycl3_V16detail24find_device_intersectionERKSt6vectorINS0_13kernel_bundleILNS0_12bundle_stateE1EEESaIS5_EE
 _ZN4sycl3_V16detail26isDeviceGlobalUsedInKernelEPKv
 _ZN4sycl3_V16detail27getPixelCoordLinearFiltModeENS0_3vecIfLi4EEENS0_15addressing_modeENS0_5rangeILi3EEERS3_
@@ -3400,7 +3398,6 @@ _ZN4sycl3_V16detail2pi9assertionEbPKc
 _ZN4sycl3_V16detail2pi9getPluginILNS0_7backendE1EEERKSt10shared_ptrINS1_6pluginEEv
 _ZN4sycl3_V16detail2pi9getPluginILNS0_7backendE2EEERKSt10shared_ptrINS1_6pluginEEv
 _ZN4sycl3_V16detail2pi9getPluginILNS0_7backendE3EEERKSt10shared_ptrINS1_6pluginEEv
-_ZN4sycl3_V16detail2pi9getPluginILNS0_7backendE5EEERKSt10shared_ptrINS1_6pluginEEv
 _ZN4sycl3_V16detail2pi9getPluginILNS0_7backendE6EEERKSt10shared_ptrINS1_6pluginEEv
 _ZN4sycl3_V16detail30UnsampledImageAccessorBaseHost10getAccDataEv
 _ZN4sycl3_V16detail30UnsampledImageAccessorBaseHost6getPtrEv