[SYCL][ESIMD] Software-emulation preparation

buildbot/configure.py

@@ -25,6 +25,7 @@ def do_configure(args):
     llvm_enable_projects = 'clang;' + llvm_external_projects
     libclc_targets_to_build = ''
     sycl_build_pi_cuda = 'OFF'
+    sycl_build_pi_esimd_cpu = 'ON'
     sycl_werror = 'ON'
     llvm_enable_assertions = 'ON'
     llvm_enable_doxygen = 'OFF'
@@ -44,6 +45,9 @@ def do_configure(args):
         libclc_targets_to_build = 'nvptx64--;nvptx64--nvidiacl'
         sycl_build_pi_cuda = 'ON'
 
+    if args.disable_esimd_cpu:
+        sycl_build_pi_esimd_cpu = 'OFF'
+
     if args.no_werror:
         sycl_werror = 'OFF'
 
@@ -86,7 +90,8 @@ def do_configure(args):
         "-DLLVM_ENABLE_SPHINX={}".format(llvm_enable_sphinx),
         "-DBUILD_SHARED_LIBS={}".format(llvm_build_shared_libs),
         "-DSYCL_ENABLE_XPTI_TRACING={}".format(sycl_enable_xpti_tracing),
-        "-DLLVM_ENABLE_LLD={}".format(llvm_enable_lld)
+        "-DLLVM_ENABLE_LLD={}".format(llvm_enable_lld),
+        "-DSYCL_BUILD_PI_ESIMD_CPU={}".format(sycl_build_pi_esimd_cpu)
     ]
 
     if args.l0_headers and args.l0_loader:
@@ -147,6 +152,7 @@ def main():
                         metavar="BUILD_TYPE", default="Release", help="build type: Debug, Release")
     parser.add_argument("--cuda", action='store_true', help="switch from OpenCL to CUDA")
     parser.add_argument("--arm", action='store_true', help="build ARM support rather than x86")
+    parser.add_argument("--disable-esimd-cpu", action='store_true', help="build without ESIMD_CPU support")
     parser.add_argument("--no-assertions", action='store_true', help="build without assertions")
     parser.add_argument("--docs", action='store_true', help="build Doxygen documentation")
     parser.add_argument("--no-werror", action='store_true', help="Don't treat warnings as errors")

sycl/CMakeLists.txt

@@ -250,6 +250,7 @@ set( SYCL_TOOLCHAIN_DEPLOY_COMPONENTS
      sycl
      pi_opencl
      pi_level_zero
+     pi_esimd_cpu
      libsycldevice
      ${XPTIFW_LIBS}
 )

sycl/include/CL/sycl/INTEL/esimd/detail/emu/esimd_emu_functions_v1.h

@@ -0,0 +1,46 @@
+//==----- esimd_emu_functions_v1.h - DPC++ Explicit SIMD API ---------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+/// \file esimd_emu_functions_v1.h
+///
+/// \ingroup sycl_pi_esimd_cpu
+
+#pragma once
+
+// <cstdint> for 'uint32_t' type is included in upper-level device
+// interface file ('esimdcpu_device_interface.hpp')
+
+// This file defines function interfaces for ESIMD CPU Emulation
+// (ESIMD_CPU) to access LibCM CPU emulation functionalities from
+// kernel applications running under emulation
+
+// CM CPU Emulation Info :
+// https://github.com/intel/cm-cpu-emulation
+
+// Function pointers (*_ptr) with 'cm/__cm' prefixes correspond to
+// LibCM functions with same name
+// e.g.: cm_fence_ptr -> cm_fence() in LibCM
+
+// Function pointers (*_ptr) with 'sycl_' prefix correspond to LibCM
+// functions dedicated to SYCL support
+// e.g.: sycl_get_surface_base_addr_ptr
+// -> get_surface_base_addr(int) in LibCM
+
+/****** DO NOT MODIFY following function pointers ******/
+/****** No reordering, No renaming, No removal    ******/
+
+// Intrinsics
+void (*cm_barrier_ptr)(void);
+void (*cm_sbarrier_ptr)(uint32_t);
+void (*cm_fence_ptr)(void);
+
+// libcm functionalities used for intrinsics such as
+// surface/buffer/slm access
+char *(*sycl_get_surface_base_addr_ptr)(int);
+char *(*__cm_emu_get_slm_ptr)(void);
+void (*cm_slm_init_ptr)(size_t);

sycl/include/CL/sycl/INTEL/esimd/detail/emu/esimdcpu_device_interface.hpp

@@ -0,0 +1,119 @@
+//==----- esimdcpu_device_interface.hpp - DPC++ Explicit SIMD API ---------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+/// \file esimdcpu_device_interface.hpp
+/// Declarations for ESIMD_CPU-device specific definitions.
+/// ESIMD intrinsic and LibCM functionalities required by intrinsic defined
+///
+/// This interface is for ESIMD intrinsic emulation implementations
+/// such as slm_access to access ESIMD_CPU specific-support therefore
+/// it has to be defined and shared as include directory
+///
+/// \ingroup sycl_pi_esimd_cpu
+
+#pragma once
+
+#include <CL/sycl/detail/pi.hpp>
+
+// cstdint-type fields such as 'uint32_t' are to be used in funtion
+// pointer table file ('esimd_emu_functions_v1.h') included in 'struct
+// ESIMDDeviceInterface' definition.
+#include <cstdint>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {
+
+/// This is the device interface version required (and used) by this
+/// implementation of the ESIMD CPU emulator.
+#define ESIMD_DEVICE_INTERFACE_VERSION 1
+
+// 'ESIMDDeviceInterface' structure defines interface for ESIMD CPU
+// emulation (ESIMD_CPU) to access LibCM CPU emulation functionalities
+// from kernel application under emulation.
+
+// Header files included in the structure contains only function
+// pointers to access CM functionalities. Only new function can be
+// added - reordering, changing, or removing existing function pointer
+// is not allowed.
+
+// Whenever a new function(s) is added to this interface, a new header
+// file must be added following naming convention that contains
+// version number such as 'v1' from 'ESIMD_DEVICE_INTERFACE_VERSION'.
+struct ESIMDDeviceInterface {
+  uintptr_t version;
+  void *reserved;
+
+  ESIMDDeviceInterface();
+#include "esimd_emu_functions_v1.h"
+};
+
+// Denotes the data version used by the implementation.
+// Increment whenever the 'data' field interpretation within PluginOpaqueData is
+// changed.
+#define ESIMD_EMU_PLUGIN_OPAQUE_DATA_VERSION 0
+/// This structure denotes a ESIMD EMU plugin-specific data returned via the
+/// piextPluginGetOpaqueData PI call. Depending on the \c version field, the
+/// second \c data field can be interpreted differently.
+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::esimd_cpu>(nullptr);
+
+  ESIMDEmuPluginOpaqueData *OpaqueData =
+      reinterpret_cast<ESIMDEmuPluginOpaqueData *>(PIOpaqueData);
+
+  // First check if opaque data version is compatible.
+  if (OpaqueData->version != ESIMD_EMU_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_EMU_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_EMU_PLUGIN_OPAQUE_DATA_VERSION
+
+} // namespace detail
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/backend_types.hpp

@@ -23,7 +23,8 @@ enum class backend : char {
   opencl = 1,
   level_zero = 2,
   cuda = 3,
-  all = 4
+  esimd_cpu = 4,
+  all = 5
 };
 
 template <backend Backend, typename SYCLObjectT> struct interop;
@@ -51,6 +52,9 @@ inline std::ostream &operator<<(std::ostream &Out, backend be) {
   case backend::cuda:
     Out << "cuda";
     break;
+  case backend::esimd_cpu:
+    Out << "esimd_cpu";
+    break;
   case backend::all:
     Out << "all";
   }

sycl/include/CL/sycl/detail/pi.def

@@ -128,6 +128,9 @@ _PI_API(piextUSMGetMemAllocInfo)
 
 _PI_API(piextKernelSetArgMemObj)
 _PI_API(piextKernelSetArgSampler)
+
+_PI_API(piextPluginGetOpaqueData)
+
 _PI_API(piTearDown)
 
 #undef _PI_API

sycl/include/CL/sycl/detail/pi.h

@@ -1634,6 +1634,15 @@ __SYCL_EXPORT pi_result piextUSMGetMemAllocInfo(
     pi_context context, const void *ptr, pi_mem_info param_name,
     size_t param_value_size, void *param_value, size_t *param_value_size_ret);
 
+/// API to get Plugin internal data, opaque to SYCL RT. Some devices whose
+/// device code is compiled by the host compiler (e.g. CPU emulators) may use it
+/// to access some device code functionality implemented in/behind the plugin.
+/// \param opaque_data_param - unspecified argument, interpretation is specific
+/// to a plugin \param opaque_data_return - placeholder for the returned opaque
+/// data.
+__SYCL_EXPORT pi_result piextPluginGetOpaqueData(void *opaque_data_param,
+                                                 void **opaque_data_return);
+
 /// API to notify that the plugin should clean up its resources.
 /// No PI calls should be made until the next piPluginInit call.
 /// \param PluginParameter placeholder for future use, currenly not used.

sycl/include/CL/sycl/detail/pi.hpp

@@ -44,6 +44,10 @@ enum class PiApiKind {
 #include <CL/sycl/detail/pi.def>
 };
 class plugin;
+
+template <cl::sycl::backend BE>
+__SYCL_EXPORT void *getPluginOpaqueData(void *opaquedata_arg);
+
 namespace pi {
 
 // The SYCL_PI_TRACE sets what we will trace.
@@ -61,10 +65,12 @@ 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_CPU_PLUGIN_NAME "pi_esimd_cpu.dll"
 #else
 #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_CPU_PLUGIN_NAME "libpi_esimd_cpu.so"
 #endif
 
 // Report error and no return (keeps compiler happy about no return statements).
@@ -149,7 +155,7 @@ extern std::shared_ptr<plugin> GlobalPlugin;
 const vector_class<plugin> &initialize();
 
 // Get the plugin serving given backend.
-template <backend BE> const plugin &getPlugin();
+template <backend BE> __SYCL_EXPORT const plugin &getPlugin();
 
 // Utility Functions to get Function Name for a PI Api.
 template <PiApiKind PiApiOffset> struct PiFuncInfo {};

sycl/plugins/CMakeLists.txt

@@ -10,3 +10,7 @@ endif()
 
 add_subdirectory(opencl)
 add_subdirectory(level_zero)
+
+if (SYCL_BUILD_PI_ESIMD_CPU)
+  add_subdirectory(esimd_cpu)
+endif()

sycl/plugins/esimd_cpu/CMakeLists.txt

@@ -0,0 +1,46 @@
+
+# PI Esimd CPU  library
+# Create Shared library for libpi_esimd_cpu.so.
+
+include_directories("${sycl_inc_dir}")
+include_directories(${OpenCL_INCLUDE_DIR})
+include_directories(${LIBCMRT_INCLUDE})
+
+add_library(pi_esimd_cpu SHARED
+  "${sycl_inc_dir}/CL/sycl/detail/pi.h"
+  "pi_esimd_cpu.cpp"
+)
+
+if (MSVC)
+  # by defining __SYCL_BUILD_SYCL_DLL, we can use __declspec(dllexport)
+  # which are individually tagged for all pi* symbols in pi.h
+  target_compile_definitions(pi_esimd_cpu PRIVATE __SYCL_BUILD_SYCL_DLL)
+else()
+  # we set the visibility of all symbols 'hidden' by default.
+  # In pi.h file, we set exported symbols with visibility==default individually
+  target_compile_options(pi_esimd_cpu PUBLIC -fvisibility=hidden)
+
+  # This script file is used to allow exporting pi* symbols only.
+  # All other symbols are regarded as local (hidden)
+  set(linker_script "${CMAKE_CURRENT_SOURCE_DIR}/../ld-version-script.txt")
+
+  # Filter symbols based on the scope defined in the script file,
+  # and export pi* function symbols in the library.
+  target_link_libraries( pi_esimd_cpu
+    PRIVATE "-Wl,--version-script=${linker_script}"
+  )
+endif()
+
+add_dependencies(sycl-toolchain pi_esimd_cpu)
+
+add_dependencies(pi_esimd_cpu
+  OpenCL-Headers)
+
+target_link_libraries(pi_esimd_cpu PRIVATE sycl)
+set_target_properties(pi_esimd_cpu PROPERTIES LINKER_LANGUAGE CXX)
+
+add_common_options(pi_esimd_cpu)
+
+install(TARGETS pi_esimd_cpu
+    LIBRARY DESTINATION "lib${LLVM_LIBDIR_SUFFIX}" COMPONENT pi_esimd_cpu
+    RUNTIME DESTINATION "bin" COMPONENT pi_esimd_cpu)