From b01084019aaa6f07eb9d5621f15dbe613cf168d1 Mon Sep 17 00:00:00 2001 From: Dongkyun Ahn Date: Fri, 19 Nov 2021 15:50:14 -0800 Subject: [PATCH 01/18] Enabling ESIMD_EMU support build by default --- buildbot/configure.py | 5 --- sycl/CMakeLists.txt | 14 +++----- sycl/plugins/CMakeLists.txt | 4 +-- .../esimd_emulator/pi_esimd_emulator.cpp | 33 ++++++++++++++----- 4 files changed, 30 insertions(+), 26 deletions(-) diff --git a/buildbot/configure.py b/buildbot/configure.py index e30b5e860c9db..6b3f1692841b1 100644 --- a/buildbot/configure.py +++ b/buildbot/configure.py @@ -29,7 +29,6 @@ def do_configure(args): libclc_targets_to_build = '' libclc_gen_remangled_variants = 'OFF' sycl_build_pi_cuda = 'OFF' - sycl_build_pi_esimd_emulator = 'OFF' sycl_build_pi_hip = 'OFF' sycl_build_pi_hip_platform = 'AMD' sycl_clang_extra_flags = '' @@ -55,9 +54,6 @@ def do_configure(args): if args.arm: llvm_targets_to_build = 'ARM;AArch64' - if args.enable_esimd_cpu_emulation: - sycl_build_pi_esimd_emulator = 'ON' - if args.cuda or args.hip: llvm_enable_projects += ';libclc' @@ -131,7 +127,6 @@ def do_configure(args): "-DBUILD_SHARED_LIBS={}".format(llvm_build_shared_libs), "-DSYCL_ENABLE_XPTI_TRACING={}".format(sycl_enable_xpti_tracing), "-DLLVM_ENABLE_LLD={}".format(llvm_enable_lld), - "-DSYCL_BUILD_PI_ESIMD_EMULATOR={}".format(sycl_build_pi_esimd_emulator), "-DXPTI_ENABLE_WERROR={}".format(xpti_enable_werror), "-DSYCL_CLANG_EXTRA_FLAGS={}".format(sycl_clang_extra_flags) ] diff --git a/sycl/CMakeLists.txt b/sycl/CMakeLists.txt index 38b6600d55244..58b9aba429427 100644 --- a/sycl/CMakeLists.txt +++ b/sycl/CMakeLists.txt @@ -332,17 +332,11 @@ if(SYCL_BUILD_PI_HIP) list(APPEND SYCL_TOOLCHAIN_DEPLOY_COMPONENTS libspirv-builtins pi_hip) 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() +list(APPEND SYCL_TOOLCHAIN_DEPLOY_COMPONENTS pi_esimd_emulator libcmrt-headers) +if (MSVC) + list(APPEND SYCL_TOOLCHAIN_DEPLOY_COMPONENTS libcmrt-libs libcmrt-dlls) 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) + list(APPEND SYCL_TOOLCHAIN_DEPLOY_COMPONENTS libcmrt-sos) endif() # Use it as fake dependency in order to force another command(s) to execute. diff --git a/sycl/plugins/CMakeLists.txt b/sycl/plugins/CMakeLists.txt index e08cf0e3b1c65..89aca08fa3565 100644 --- a/sycl/plugins/CMakeLists.txt +++ b/sycl/plugins/CMakeLists.txt @@ -18,7 +18,5 @@ add_subdirectory(level_zero) # TODO : Remove 'if (NOT MSVC)' when CM_EMU supports Windows # environment if (NOT MSVC) - if (SYCL_BUILD_PI_ESIMD_EMULATOR) - add_subdirectory(esimd_emulator) - endif() + add_subdirectory(esimd_emulator) endif() diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index fa6c335458a70..89554697f2dbf 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -123,6 +123,9 @@ static sycl::detail::ESIMDEmuPluginOpaqueData *PiESimdDeviceAccess; // interface header file #define ESIMDEmuPluginInterfaceVersion 1 +// For PI_DEVICE_INFO_DRIVER_VERSION info +static char ESimdEmuVersionString[32]; + using IDBuilder = sycl::detail::Builder; template @@ -398,7 +401,7 @@ pi_result piPlatformGetInfo(pi_platform Platform, pi_platform_info ParamName, return ReturnValue("Intel(R) Corporation"); case PI_PLATFORM_INFO_VERSION: - return ReturnValue(Platform->CmEmuVersion); + return ReturnValue(Platform->CmEmuVersion.c_str()); case PI_PLATFORM_INFO_PROFILE: return ReturnValue("FULL_PROFILE"); @@ -499,7 +502,13 @@ pi_result piDeviceGetInfo(pi_device Device, pi_device_info ParamName, case PI_DEVICE_INFO_IMAGE_SUPPORT: return ReturnValue(pi_bool{true}); case PI_DEVICE_INFO_DRIVER_VERSION: - return ReturnValue("0.0.1"); + /// Combination of ESIMDEmuPluginDataVersion and + /// ESIMDEmuPluginInterfaceVersion : 0.a.b + /// a : ESIMDEmuPluginInterfaceVersion + /// b : ESIMDEmuPluginDataVersion + sprintf(ESimdEmuVersionString, "0.%d.%d", ESIMDEmuPluginInterfaceVersion, + ESIMDEmuPluginDataVersion); + return ReturnValue(ESimdEmuVersionString); case PI_DEVICE_INFO_VENDOR: return ReturnValue("Intel(R) Corporation"); case PI_DEVICE_INFO_IMAGE2D_MAX_WIDTH: @@ -513,6 +522,20 @@ pi_result piDeviceGetInfo(pi_device Device, pi_device_info ParamName, // cl_khr_fp64, cl_khr_int64_base_atomics, // cl_khr_int64_extended_atomics return ReturnValue(""); + case PI_DEVICE_INFO_VERSION: + // CM_EMU release version from + // https://github.com/intel/cm-cpu-emulation/releases + return ReturnValue("1.0.7-CM_EMU"); + case PI_DEVICE_INFO_COMPILER_AVAILABLE: + return ReturnValue(pi_bool{false}); + case PI_DEVICE_INFO_LINKER_AVAILABLE: + return ReturnValue(pi_bool{false}); + case PI_DEVICE_INFO_MAX_COMPUTE_UNITS: + return ReturnValue(pi_uint32{256}); + case PI_DEVICE_INFO_PARTITION_MAX_SUB_DEVICES: + return ReturnValue(pi_uint32{0}); + case PI_DEVICE_INFO_PARTITION_PROPERTIES: + return ReturnValue(pi_device_partition_property{0}); #define UNSUPPORTED_INFO(info) \ case info: \ @@ -523,9 +546,6 @@ pi_result piDeviceGetInfo(pi_device Device, pi_device_info ParamName, break; UNSUPPORTED_INFO(PI_DEVICE_INFO_VENDOR_ID) - UNSUPPORTED_INFO(PI_DEVICE_INFO_COMPILER_AVAILABLE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_LINKER_AVAILABLE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_MAX_COMPUTE_UNITS) UNSUPPORTED_INFO(PI_DEVICE_INFO_MAX_WORK_ITEM_DIMENSIONS) UNSUPPORTED_INFO(PI_DEVICE_INFO_MAX_WORK_GROUP_SIZE) UNSUPPORTED_INFO(PI_DEVICE_INFO_MAX_WORK_ITEM_SIZES) @@ -535,10 +555,7 @@ pi_result piDeviceGetInfo(pi_device Device, pi_device_info ParamName, UNSUPPORTED_INFO(PI_DEVICE_INFO_GLOBAL_MEM_SIZE) UNSUPPORTED_INFO(PI_DEVICE_INFO_LOCAL_MEM_SIZE) UNSUPPORTED_INFO(PI_DEVICE_INFO_AVAILABLE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_VERSION) - UNSUPPORTED_INFO(PI_DEVICE_INFO_PARTITION_MAX_SUB_DEVICES) UNSUPPORTED_INFO(PI_DEVICE_INFO_REFERENCE_COUNT) - UNSUPPORTED_INFO(PI_DEVICE_INFO_PARTITION_PROPERTIES) UNSUPPORTED_INFO(PI_DEVICE_INFO_PARTITION_AFFINITY_DOMAIN) UNSUPPORTED_INFO(PI_DEVICE_INFO_PARTITION_TYPE) UNSUPPORTED_INFO(PI_DEVICE_INFO_OPENCL_C_VERSION) From f8ae218e0099c5e853334b55ae3074e8d8e0fd5d Mon Sep 17 00:00:00 2001 From: Dongkyun Ahn Date: Tue, 30 Nov 2021 11:46:22 -0800 Subject: [PATCH 02/18] Replacing ESIMDEmu with ESIMDCPU for legacy CM_EMU - To be replaced at CM_EMU update --- sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 89554697f2dbf..34f3f4ee79099 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -28,7 +28,8 @@ #include #include -#include +// TODO : Rename esimdcpu to esimdemu for next CM_EMU release +#include #include #include @@ -245,7 +246,9 @@ template class libCMBatch { GroupDim[I] = (uint32_t)(GlobalSize[I] / LocalSize[I]); } - EsimdemuKernel Esimdemu((fptrVoid)InvokeLambda, GroupDim, SpaceDim); + // TODO : Replace ESimdCPUKernel with EsimdemuKernel for next + // CM_EMU release + ESimdCPUKernel Esimdemu((fptrVoid)InvokeLambda, GroupDim, SpaceDim); Esimdemu.launchMT(sizeof(struct LambdaWrapper), WrappedLambda.get()); } From 1d8cb9c583cc4bb7db7db4f8f05a735deba35836 Mon Sep 17 00:00:00 2001 From: Dongkyun Ahn Date: Wed, 1 Dec 2021 14:44:57 -0800 Subject: [PATCH 03/18] PI Device info update --- .../esimd_emulator/pi_esimd_emulator.cpp | 204 ++++++++++++------ 1 file changed, 143 insertions(+), 61 deletions(-) diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 34f3f4ee79099..41e6a0a873eeb 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -539,6 +539,149 @@ pi_result piDeviceGetInfo(pi_device Device, pi_device_info ParamName, return ReturnValue(pi_uint32{0}); case PI_DEVICE_INFO_PARTITION_PROPERTIES: return ReturnValue(pi_device_partition_property{0}); + case PI_DEVICE_INFO_VENDOR_ID: + // '0x8086' : 'Intel HD graphics vendor ID' + return ReturnValue(pi_uint32{0x8086}); + case PI_DEVICE_INFO_LOCAL_MEM_SIZE: + // Default SLM_MAX_SIZE from CM_EMU + return ReturnValue(pi_uint32{65536}); + case PI_DEVICE_INFO_MAX_WORK_GROUP_SIZE: + return ReturnValue(size_t{256}); + case PI_DEVICE_INFO_MEM_BASE_ADDR_ALIGN: + // Imported from level_zero + return ReturnValue(pi_uint32{8}); + case PI_DEVICE_INFO_IMAGE3D_MAX_WIDTH: + case PI_DEVICE_INFO_IMAGE3D_MAX_HEIGHT: + case PI_DEVICE_INFO_IMAGE3D_MAX_DEPTH: + // Default minimum values required by the SYCL specification. + return ReturnValue(size_t{2048}); + case PI_DEVICE_INFO_MAX_WORK_ITEM_DIMENSIONS: + return ReturnValue(pi_uint32{3}); + case PI_DEVICE_INFO_PARTITION_TYPE: + return ReturnValue(pi_device_partition_property{0}); + case PI_DEVICE_INFO_OPENCL_C_VERSION: + return ReturnValue(""); + case PI_DEVICE_INFO_QUEUE_PROPERTIES: + return ReturnValue(pi_queue_properties{PI_QUEUE_ON_DEVICE}); + case PI_DEVICE_INFO_MAX_WORK_ITEM_SIZES: { + struct { + size_t Arr[3]; + } MaxGroupSize = {{256, 256, 1}}; + return ReturnValue(MaxGroupSize); + } + case PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_CHAR: + case PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_SHORT: + case PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_INT: + case PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_LONG: + case PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_FLOAT: + case PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_DOUBLE: + case PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_HALF: + case PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_CHAR: + case PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_SHORT: + case PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_INT: + case PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_LONG: + case PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_FLOAT: + case PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_DOUBLE: + case PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_HALF: + return ReturnValue(pi_uint32{1}); + + // Imported from level_zero + case PI_DEVICE_INFO_USM_HOST_SUPPORT: + case PI_DEVICE_INFO_USM_DEVICE_SUPPORT: + case PI_DEVICE_INFO_USM_SINGLE_SHARED_SUPPORT: + case PI_DEVICE_INFO_USM_CROSS_SHARED_SUPPORT: + case PI_DEVICE_INFO_USM_SYSTEM_SHARED_SUPPORT: { + pi_uint64 Supported = 0; + // TODO[1.0]: how to query for USM support now? + if (true) { + // TODO: Use ze_memory_access_capabilities_t + Supported = PI_USM_ACCESS | PI_USM_ATOMIC_ACCESS | + PI_USM_CONCURRENT_ACCESS | PI_USM_CONCURRENT_ATOMIC_ACCESS; + } + return ReturnValue(Supported); + } + case PI_DEVICE_INFO_ADDRESS_BITS: + return ReturnValue( + pi_uint32{sizeof(void *) * std::numeric_limits::digits}); + case PI_DEVICE_INFO_MAX_CLOCK_FREQUENCY: + return ReturnValue(pi_uint32{1000}); + case PI_DEVICE_INFO_ENDIAN_LITTLE: + return ReturnValue(pi_bool{true}); + case PI_DEVICE_INFO_AVAILABLE: + return ReturnValue(pi_bool{true}); + case PI_DEVICE_INFO_MAX_READ_IMAGE_ARGS: + case PI_DEVICE_INFO_MAX_WRITE_IMAGE_ARGS: + /// TODO : Check + return ReturnValue(pi_uint32{0}); + case PI_DEVICE_INFO_IMAGE_MAX_BUFFER_SIZE: + /// TODO : Check. CM_MAX_1D_SURF_WIDTH from CM_EMU + return ReturnValue(size_t{0x80000000}); + case PI_DEVICE_INFO_IMAGE_MAX_ARRAY_SIZE: + /// TODO : Check + return ReturnValue(size_t{0}); + case PI_DEVICE_INFO_MAX_SAMPLERS: + /// TODO : Check. CM_MAX_SAMPLERS_PER_KERNEL from CM_EMU + return ReturnValue(pi_uint32{16}); + case PI_DEVICE_INFO_MAX_PARAMETER_SIZE: + /// TODO : Check + return ReturnValue(size_t{32}); + case PI_DEVICE_INFO_HALF_FP_CONFIG: + case PI_DEVICE_INFO_SINGLE_FP_CONFIG: + case PI_DEVICE_INFO_DOUBLE_FP_CONFIG: { + /// TODO : Check. half_type.hpp from CM_EMU + uint64_t FPValue = PI_FP_CORRECTLY_ROUNDED_DIVIDE_SQRT | + PI_FP_ROUND_TO_NEAREST | PI_FP_ROUND_TO_ZERO | + PI_FP_ROUND_TO_INF | PI_FP_INF_NAN | PI_FP_DENORM | + PI_FP_FMA; + return ReturnValue(pi_uint64{FPValue}); + } + case PI_DEVICE_INFO_GLOBAL_MEM_CACHE_TYPE: + return ReturnValue(PI_DEVICE_MEM_CACHE_TYPE_READ_WRITE_CACHE); + case PI_DEVICE_INFO_GLOBAL_MEM_CACHELINE_SIZE: + // TODO : CHECK + return ReturnValue(pi_uint32{64}); + case PI_DEVICE_INFO_GLOBAL_MEM_CACHE_SIZE: + // TODO : CHECK + return ReturnValue(pi_uint64{0}); + case PI_DEVICE_INFO_GLOBAL_MEM_SIZE: + // TODO : CHECK + return ReturnValue(pi_uint64{0}); + case PI_DEVICE_INFO_MAX_CONSTANT_BUFFER_SIZE: + // TODO : CHECK + return ReturnValue(pi_uint64{0}); + case PI_DEVICE_INFO_MAX_CONSTANT_ARGS: + // TODO : CHECK + return ReturnValue(pi_uint32{64}); + case PI_DEVICE_INFO_LOCAL_MEM_TYPE: + // TODO : CHECK + return ReturnValue(PI_DEVICE_LOCAL_MEM_TYPE_LOCAL); + case PI_DEVICE_INFO_ERROR_CORRECTION_SUPPORT: + return ReturnValue(pi_bool{false}); + case PI_DEVICE_INFO_PROFILING_TIMER_RESOLUTION: + // TODO : CHECK + return ReturnValue(size_t{0}); + case PI_DEVICE_INFO_BUILT_IN_KERNELS: + // TODO : CHECK + return ReturnValue(""); + case PI_DEVICE_INFO_PRINTF_BUFFER_SIZE: + // TODO : CHECK + return ReturnValue(size_t{1024}); + case PI_DEVICE_INFO_PREFERRED_INTEROP_USER_SYNC: + return ReturnValue(pi_bool{false}); + case PI_DEVICE_INFO_PARTITION_AFFINITY_DOMAIN: + return ReturnValue(pi_device_affinity_domain{0}); + case PI_DEVICE_INFO_MAX_MEM_ALLOC_SIZE: + // TODO : CHECK + return ReturnValue(pi_uint64{0}); + case PI_DEVICE_INFO_EXECUTION_CAPABILITIES: + // TODO : CHECK + return ReturnValue( + pi_device_exec_capabilities{PI_DEVICE_EXEC_CAPABILITIES_KERNEL}); + case PI_DEVICE_INFO_PROFILE: + return ReturnValue("FULL_PROFILE"); + case PI_DEVICE_INFO_REFERENCE_COUNT: + // TODO : CHECK + return ReturnValue(pi_uint32{0}); #define UNSUPPORTED_INFO(info) \ case info: \ @@ -548,71 +691,10 @@ pi_result piDeviceGetInfo(pi_device Device, pi_device_info ParamName, DIE_NO_IMPLEMENTATION; \ break; - UNSUPPORTED_INFO(PI_DEVICE_INFO_VENDOR_ID) - UNSUPPORTED_INFO(PI_DEVICE_INFO_MAX_WORK_ITEM_DIMENSIONS) - UNSUPPORTED_INFO(PI_DEVICE_INFO_MAX_WORK_GROUP_SIZE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_MAX_WORK_ITEM_SIZES) - UNSUPPORTED_INFO(PI_DEVICE_INFO_MAX_CLOCK_FREQUENCY) - UNSUPPORTED_INFO(PI_DEVICE_INFO_ADDRESS_BITS) - UNSUPPORTED_INFO(PI_DEVICE_INFO_MAX_MEM_ALLOC_SIZE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_GLOBAL_MEM_SIZE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_LOCAL_MEM_SIZE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_AVAILABLE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_REFERENCE_COUNT) - UNSUPPORTED_INFO(PI_DEVICE_INFO_PARTITION_AFFINITY_DOMAIN) - UNSUPPORTED_INFO(PI_DEVICE_INFO_PARTITION_TYPE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_OPENCL_C_VERSION) - UNSUPPORTED_INFO(PI_DEVICE_INFO_PREFERRED_INTEROP_USER_SYNC) - UNSUPPORTED_INFO(PI_DEVICE_INFO_PRINTF_BUFFER_SIZE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_PROFILE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_BUILT_IN_KERNELS) - UNSUPPORTED_INFO(PI_DEVICE_INFO_QUEUE_PROPERTIES) - UNSUPPORTED_INFO(PI_DEVICE_INFO_EXECUTION_CAPABILITIES) - UNSUPPORTED_INFO(PI_DEVICE_INFO_ENDIAN_LITTLE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_ERROR_CORRECTION_SUPPORT) - UNSUPPORTED_INFO(PI_DEVICE_INFO_PROFILING_TIMER_RESOLUTION) - UNSUPPORTED_INFO(PI_DEVICE_INFO_LOCAL_MEM_TYPE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_MAX_CONSTANT_ARGS) - UNSUPPORTED_INFO(PI_DEVICE_INFO_MAX_CONSTANT_BUFFER_SIZE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_GLOBAL_MEM_CACHE_TYPE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_GLOBAL_MEM_CACHELINE_SIZE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_GLOBAL_MEM_CACHE_SIZE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_MAX_PARAMETER_SIZE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_MEM_BASE_ADDR_ALIGN) - UNSUPPORTED_INFO(PI_DEVICE_INFO_MAX_SAMPLERS) - UNSUPPORTED_INFO(PI_DEVICE_INFO_MAX_READ_IMAGE_ARGS) - UNSUPPORTED_INFO(PI_DEVICE_INFO_MAX_WRITE_IMAGE_ARGS) - UNSUPPORTED_INFO(PI_DEVICE_INFO_SINGLE_FP_CONFIG) - UNSUPPORTED_INFO(PI_DEVICE_INFO_HALF_FP_CONFIG) - UNSUPPORTED_INFO(PI_DEVICE_INFO_DOUBLE_FP_CONFIG) - UNSUPPORTED_INFO(PI_DEVICE_INFO_IMAGE3D_MAX_WIDTH) - UNSUPPORTED_INFO(PI_DEVICE_INFO_IMAGE3D_MAX_HEIGHT) - UNSUPPORTED_INFO(PI_DEVICE_INFO_IMAGE3D_MAX_DEPTH) - UNSUPPORTED_INFO(PI_DEVICE_INFO_IMAGE_MAX_BUFFER_SIZE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_IMAGE_MAX_ARRAY_SIZE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_CHAR) - UNSUPPORTED_INFO(PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_CHAR) - UNSUPPORTED_INFO(PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_SHORT) - UNSUPPORTED_INFO(PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_SHORT) - UNSUPPORTED_INFO(PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_INT) - UNSUPPORTED_INFO(PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_INT) - UNSUPPORTED_INFO(PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_LONG) - UNSUPPORTED_INFO(PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_LONG) - UNSUPPORTED_INFO(PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_FLOAT) - UNSUPPORTED_INFO(PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_FLOAT) - UNSUPPORTED_INFO(PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_DOUBLE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_DOUBLE) - UNSUPPORTED_INFO(PI_DEVICE_INFO_NATIVE_VECTOR_WIDTH_HALF) - UNSUPPORTED_INFO(PI_DEVICE_INFO_PREFERRED_VECTOR_WIDTH_HALF) UNSUPPORTED_INFO(PI_DEVICE_INFO_MAX_NUM_SUB_GROUPS) UNSUPPORTED_INFO(PI_DEVICE_INFO_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS) UNSUPPORTED_INFO(PI_DEVICE_INFO_SUB_GROUP_SIZES_INTEL) UNSUPPORTED_INFO(PI_DEVICE_INFO_IL_VERSION) - UNSUPPORTED_INFO(PI_DEVICE_INFO_USM_HOST_SUPPORT) - UNSUPPORTED_INFO(PI_DEVICE_INFO_USM_DEVICE_SUPPORT) - UNSUPPORTED_INFO(PI_DEVICE_INFO_USM_SINGLE_SHARED_SUPPORT) - UNSUPPORTED_INFO(PI_DEVICE_INFO_USM_CROSS_SHARED_SUPPORT) - UNSUPPORTED_INFO(PI_DEVICE_INFO_USM_SYSTEM_SHARED_SUPPORT) #undef UNSUPPORTED_INFO default: From 183a2b326367df160aaf5fd9ac55d88689edfbb7 Mon Sep 17 00:00:00 2001 From: Dongkyun Ahn Date: Wed, 1 Dec 2021 22:37:05 -0800 Subject: [PATCH 04/18] Removing command line option for esimd_cpu_emulation --- buildbot/configure.py | 1 - 1 file changed, 1 deletion(-) diff --git a/buildbot/configure.py b/buildbot/configure.py index 6b3f1692841b1..3f1972fc8fa6b 100644 --- a/buildbot/configure.py +++ b/buildbot/configure.py @@ -192,7 +192,6 @@ def main(): parser.add_argument("--hip-platform", type=str, choices=['AMD', 'NVIDIA'], default='AMD', help="choose hardware platform for HIP backend") parser.add_argument("--hip-amd-arch", type=str, help="Sets AMD gpu architecture for llvm lit tests, this is only needed for the HIP backend and AMD platform") parser.add_argument("--arm", action='store_true', help="build ARM support rather than x86") - parser.add_argument("--enable-esimd-cpu-emulation", action='store_true', help="build with ESIMD_CPU emulation 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") From 340e5dd5049db0ba6165a7efabe7af04907a4203 Mon Sep 17 00:00:00 2001 From: Dongkyun Ahn Date: Thu, 2 Dec 2021 12:33:27 -0800 Subject: [PATCH 05/18] Applying reducing overhead on command creation for ESIMD_EMULATOR BE - Reducing overhead on command creation in specific case from PR#4841 - Same kernel launching flow for ESIMD_EMULATOR is in ExecCGCommand::enqueueImp() in commands.cpp --- sycl/source/handler.cpp | 15 ++++++++++++--- 1 file changed, 12 insertions(+), 3 deletions(-) diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index e1c27d2650898..9e2f23e78b2d8 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -11,6 +11,7 @@ #include #include #include +#include #include #include #include @@ -206,9 +207,17 @@ event handler::finalize() { if (MQueue->is_host()) { MHostKernel->call(MNDRDesc, NewEvent->getHostProfilingInfo()); } else { - Res = enqueueImpKernel(MQueue, MNDRDesc, MArgs, KernelBundleImpPtr, - MKernel, MKernelName, MOSModuleHandle, RawEvents, - NewEvent, nullptr); + if (MQueue->getPlugin().getBackend() == + backend::ext_intel_esimd_emulator) { + MQueue->getPlugin().call( + nullptr, reinterpret_cast(MHostKernel->getPtr()), + MNDRDesc.Dims, &MNDRDesc.GlobalOffset[0], &MNDRDesc.GlobalSize[0], + &MNDRDesc.LocalSize[0], 0, nullptr, nullptr); + } else { + Res = enqueueImpKernel(MQueue, MNDRDesc, MArgs, KernelBundleImpPtr, + MKernel, MKernelName, MOSModuleHandle, RawEvents, + NewEvent, nullptr); + } } if (CL_SUCCESS != Res) From 8ba735d0e9e4b55a8ba5ba2be6cdb7dcceb7bb03 Mon Sep 17 00:00:00 2001 From: Dongkyun Ahn Date: Mon, 6 Dec 2021 08:07:03 -0800 Subject: [PATCH 06/18] ESIMD_EMULATOR plug-in update for bringing-up intel/llvm-test-suite --- .../esimd_emulator/pi_esimd_emulator.cpp | 84 +++++++++++++++---- sycl/source/detail/scheduler/commands.cpp | 4 +- sycl/source/handler.cpp | 6 +- 3 files changed, 75 insertions(+), 19 deletions(-) diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 41e6a0a873eeb..9eb2f4a7810d8 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -360,6 +360,16 @@ extern "C" { } \ return PI_SUCCESS; +#define CASE_PI_UNSUPPORTED(not_supported) \ + case not_supported: \ + if (PrintPiTrace) { \ + std::cerr << std::endl \ + << "Unsupported PI case : " << #not_supported << " in " \ + << __FUNCTION__ << ":" << __LINE__ << "(" << __FILE__ << ")" \ + << std::endl; \ + } \ + return PI_INVALID_OPERATION; + pi_result piPlatformsGet(pi_uint32 NumEntries, pi_platform *Platforms, pi_uint32 *NumPlatforms) { @@ -528,7 +538,7 @@ pi_result piDeviceGetInfo(pi_device Device, pi_device_info ParamName, case PI_DEVICE_INFO_VERSION: // CM_EMU release version from // https://github.com/intel/cm-cpu-emulation/releases - return ReturnValue("1.0.7-CM_EMU"); + return ReturnValue("1.0"); case PI_DEVICE_INFO_COMPILER_AVAILABLE: return ReturnValue(pi_bool{false}); case PI_DEVICE_INFO_LINKER_AVAILABLE: @@ -683,20 +693,27 @@ pi_result piDeviceGetInfo(pi_device Device, pi_device_info ParamName, // TODO : CHECK return ReturnValue(pi_uint32{0}); -#define UNSUPPORTED_INFO(info) \ - case info: \ - std::cerr << std::endl \ - << "Unsupported device info = " << #info \ - << " from ESIMD_EMULATOR" << std::endl; \ - DIE_NO_IMPLEMENTATION; \ - break; - - UNSUPPORTED_INFO(PI_DEVICE_INFO_MAX_NUM_SUB_GROUPS) - UNSUPPORTED_INFO(PI_DEVICE_INFO_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS) - UNSUPPORTED_INFO(PI_DEVICE_INFO_SUB_GROUP_SIZES_INTEL) - UNSUPPORTED_INFO(PI_DEVICE_INFO_IL_VERSION) + CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_MAX_NUM_SUB_GROUPS) + CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS) + CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_SUB_GROUP_SIZES_INTEL) + CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_IL_VERSION) + + // Intel-specific extensions + CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_PCI_ADDRESS) + CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_GPU_EU_COUNT) + CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_GPU_EU_SIMD_WIDTH) + CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_GPU_SLICES) + CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_GPU_SUBSLICES_PER_SLICE) + CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_GPU_EU_COUNT_PER_SUBSLICE) + CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_MAX_MEM_BANDWIDTH) + CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_IMAGE_SRGB) + CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_ATOMIC_64) + CASE_PI_UNSUPPORTED(PI_DEVICE_INFO_ATOMIC_MEMORY_ORDER_CAPABILITIES) + CASE_PI_UNSUPPORTED(PI_EXT_ONEAPI_DEVICE_INFO_MAX_GLOBAL_WORK_GROUPS) + CASE_PI_UNSUPPORTED(PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_1D) + CASE_PI_UNSUPPORTED(PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_2D) + CASE_PI_UNSUPPORTED(PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_3D) -#undef UNSUPPORTED_INFO default: DIE_NO_IMPLEMENTATION; } @@ -999,6 +1016,13 @@ pi_result piMemImageCreate(pi_context Context, pi_mem_flags Flags, switch (ImageDesc->image_type) { case PI_MEM_TYPE_IMAGE2D: break; + + CASE_PI_UNSUPPORTED(PI_MEM_TYPE_IMAGE3D) + CASE_PI_UNSUPPORTED(PI_MEM_TYPE_IMAGE2D_ARRAY) + CASE_PI_UNSUPPORTED(PI_MEM_TYPE_IMAGE1D) + CASE_PI_UNSUPPORTED(PI_MEM_TYPE_IMAGE1D_ARRAY) + CASE_PI_UNSUPPORTED(PI_MEM_TYPE_IMAGE1D_BUFFER) + default: return PI_INVALID_MEM_OBJECT; } @@ -1012,6 +1036,18 @@ pi_result piMemImageCreate(pi_context Context, pi_mem_flags Flags, case PI_IMAGE_CHANNEL_TYPE_UNORM_INT8: BytesPerPixel = 4; break; + CASE_PI_UNSUPPORTED(PI_IMAGE_CHANNEL_TYPE_SNORM_INT8) + CASE_PI_UNSUPPORTED(PI_IMAGE_CHANNEL_TYPE_SNORM_INT16) + CASE_PI_UNSUPPORTED(PI_IMAGE_CHANNEL_TYPE_UNORM_INT16) + CASE_PI_UNSUPPORTED(PI_IMAGE_CHANNEL_TYPE_UNORM_SHORT_565) + CASE_PI_UNSUPPORTED(PI_IMAGE_CHANNEL_TYPE_UNORM_SHORT_555) + CASE_PI_UNSUPPORTED(PI_IMAGE_CHANNEL_TYPE_UNORM_INT_101010) + CASE_PI_UNSUPPORTED(PI_IMAGE_CHANNEL_TYPE_SIGNED_INT8) + CASE_PI_UNSUPPORTED(PI_IMAGE_CHANNEL_TYPE_SIGNED_INT16) + CASE_PI_UNSUPPORTED(PI_IMAGE_CHANNEL_TYPE_SIGNED_INT32) + CASE_PI_UNSUPPORTED(PI_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16) + CASE_PI_UNSUPPORTED(PI_IMAGE_CHANNEL_TYPE_HALF_FLOAT) + CASE_PI_UNSUPPORTED(PI_IMAGE_CHANNEL_TYPE_FLOAT) default: return PI_IMAGE_FORMAT_NOT_SUPPORTED; } @@ -1324,6 +1360,7 @@ pi_result piEnqueueMemBufferReadRect(pi_queue, pi_mem, pi_bool, pi_result piEnqueueMemBufferWrite(pi_queue, pi_mem, pi_bool, size_t, size_t, const void *, pi_uint32, const pi_event *, pi_event *) { + // TODO : intel/llvm_test_suite DIE_NO_IMPLEMENTATION; } @@ -1338,6 +1375,7 @@ pi_result piEnqueueMemBufferWriteRect(pi_queue, pi_mem, pi_bool, pi_result piEnqueueMemBufferCopy(pi_queue, pi_mem, pi_mem, size_t, size_t, size_t, pi_uint32, const pi_event *, pi_event *) { + // TODO : intel/llvm_test_suite DIE_NO_IMPLEMENTATION; } @@ -1346,18 +1384,21 @@ pi_result piEnqueueMemBufferCopyRect(pi_queue, pi_mem, pi_mem, pi_buff_rect_region, size_t, size_t, size_t, size_t, pi_uint32, const pi_event *, pi_event *) { + // TODO : intel/llvm_test_suite DIE_NO_IMPLEMENTATION; } pi_result piEnqueueMemBufferFill(pi_queue, pi_mem, const void *, size_t, size_t, size_t, pi_uint32, const pi_event *, pi_event *) { + // TODO : intel/llvm_test_suite DIE_NO_IMPLEMENTATION; } pi_result piEnqueueMemBufferMap(pi_queue, pi_mem, pi_bool, pi_map_flags, size_t, size_t, pi_uint32, const pi_event *, pi_event *, void **) { + // TODO : intel/llvm_test_suite DIE_NO_IMPLEMENTATION; } @@ -1439,7 +1480,9 @@ piEnqueueKernelLaunch(pi_queue Queue, pi_kernel Kernel, pi_uint32 WorkDim, return PI_INVALID_KERNEL; } - if ((WorkDim > 3) || (WorkDim == 0)) { + // WorkDim == 0 is reserved for 'single_task()' kernel with no + // argument + if (WorkDim > 3) { return PI_INVALID_WORK_GROUP_SIZE; } @@ -1461,6 +1504,12 @@ piEnqueueKernelLaunch(pi_queue Queue, pi_kernel Kernel, pi_uint32 WorkDim, } switch (WorkDim) { + case 0: + // TODO : intel/llvm_test_suite + // single_task() support - void(*)(void) + DIE_NO_IMPLEMENTATION; + break; + case 1: InvokeImpl<1>::invoke(Kernel, GlobalWorkOffset, GlobalWorkSize, LocalWorkSize); @@ -1576,6 +1625,7 @@ pi_result piextKernelSetArgPointer(pi_kernel, pi_uint32, size_t, const void *) { pi_result piextUSMEnqueueMemset(pi_queue, void *, pi_int32, size_t, pi_uint32, const pi_event *, pi_event *) { + // TODO : intel/llvm_test_suite DIE_NO_IMPLEMENTATION; } @@ -1586,6 +1636,7 @@ pi_result piextUSMEnqueueMemcpy(pi_queue, pi_bool, void *, const void *, size_t, pi_result piextUSMEnqueueMemAdvise(pi_queue, const void *, size_t, pi_mem_advice, pi_event *) { + // TODO : intel/llvm_test_suite DIE_NO_IMPLEMENTATION; } @@ -1623,7 +1674,8 @@ pi_result piextDeviceSelectBinary(pi_device, pi_device_binary *, pi_result piextUSMEnqueuePrefetch(pi_queue, const void *, size_t, pi_usm_migration_flags, pi_uint32, const pi_event *, pi_event *) { - DIE_NO_IMPLEMENTATION; + // NOP for prefetch + return PI_SUCCESS; } pi_result piextPluginGetOpaqueData(void *, void **OpaqueDataReturn) { diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index f74b85e229942..07482776bdf91 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -2157,10 +2157,12 @@ cl_int ExecCGCommand::enqueueImp() { } else { assert(MQueue->getPlugin().getBackend() == backend::ext_intel_esimd_emulator); + // Dims==0 for 'single_task() - void(void) type' + uint32_t Dims = (Args.size() > 0) ? NDRDesc.Dims : 0; MQueue->getPlugin().call( nullptr, reinterpret_cast(ExecKernel->MHostKernel->getPtr()), - NDRDesc.Dims, &NDRDesc.GlobalOffset[0], &NDRDesc.GlobalSize[0], + Dims, &NDRDesc.GlobalOffset[0], &NDRDesc.GlobalSize[0], &NDRDesc.LocalSize[0], 0, nullptr, nullptr); } diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 9e2f23e78b2d8..7970546ee7195 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -209,9 +209,11 @@ event handler::finalize() { } else { if (MQueue->getPlugin().getBackend() == backend::ext_intel_esimd_emulator) { + // Dims==0 for 'single_task() - void(void) type' + uint32_t Dims = (MArgs.size() > 0) ? MNDRDesc.Dims : 0; MQueue->getPlugin().call( - nullptr, reinterpret_cast(MHostKernel->getPtr()), - MNDRDesc.Dims, &MNDRDesc.GlobalOffset[0], &MNDRDesc.GlobalSize[0], + nullptr, reinterpret_cast(MHostKernel->getPtr()), Dims, + &MNDRDesc.GlobalOffset[0], &MNDRDesc.GlobalSize[0], &MNDRDesc.LocalSize[0], 0, nullptr, nullptr); } else { Res = enqueueImpKernel(MQueue, MNDRDesc, MArgs, KernelBundleImpPtr, From 0b9c1e7ee770572b03e14894234b0490aa8ea2c6 Mon Sep 17 00:00:00 2001 From: Dongkyun Ahn Date: Mon, 13 Dec 2021 16:08:27 -0800 Subject: [PATCH 07/18] Removing 'TODO' comments for tests from intel/llvm-test-suite - As esimd_emulator is only for ESIMD kernels, PI_APIs causing failures for non-ESIMD kernels are not going to be implemented. 'TODO' comments are removed for such PI_APIs - CMakeLists change : installtion path fix for headers imported from CM --- sycl/plugins/esimd_emulator/CMakeLists.txt | 2 +- sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp | 10 +--------- 2 files changed, 2 insertions(+), 10 deletions(-) diff --git a/sycl/plugins/esimd_emulator/CMakeLists.txt b/sycl/plugins/esimd_emulator/CMakeLists.txt index 4ce9e1ff53e61..98963faccefca 100755 --- a/sycl/plugins/esimd_emulator/CMakeLists.txt +++ b/sycl/plugins/esimd_emulator/CMakeLists.txt @@ -138,7 +138,7 @@ install(TARGETS pi_esimd_emulator # Copy CM Header files to $(INSTALL)/include/sycl/CL/ install(DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}/cm-emu_install/include/libcm/cm/ - DESTINATION ${SYCL_INCLUDE_DIR}/CL + DESTINATION ${SYCL_INCLUDE_DIR}/sycl/CL COMPONENT libcmrt-headers FILES_MATCHING PATTERN "*.h" ) diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 9eb2f4a7810d8..4fc717cbaa6e8 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -1360,7 +1360,6 @@ pi_result piEnqueueMemBufferReadRect(pi_queue, pi_mem, pi_bool, pi_result piEnqueueMemBufferWrite(pi_queue, pi_mem, pi_bool, size_t, size_t, const void *, pi_uint32, const pi_event *, pi_event *) { - // TODO : intel/llvm_test_suite DIE_NO_IMPLEMENTATION; } @@ -1375,7 +1374,6 @@ pi_result piEnqueueMemBufferWriteRect(pi_queue, pi_mem, pi_bool, pi_result piEnqueueMemBufferCopy(pi_queue, pi_mem, pi_mem, size_t, size_t, size_t, pi_uint32, const pi_event *, pi_event *) { - // TODO : intel/llvm_test_suite DIE_NO_IMPLEMENTATION; } @@ -1384,21 +1382,18 @@ pi_result piEnqueueMemBufferCopyRect(pi_queue, pi_mem, pi_mem, pi_buff_rect_region, size_t, size_t, size_t, size_t, pi_uint32, const pi_event *, pi_event *) { - // TODO : intel/llvm_test_suite DIE_NO_IMPLEMENTATION; } pi_result piEnqueueMemBufferFill(pi_queue, pi_mem, const void *, size_t, size_t, size_t, pi_uint32, const pi_event *, pi_event *) { - // TODO : intel/llvm_test_suite DIE_NO_IMPLEMENTATION; } pi_result piEnqueueMemBufferMap(pi_queue, pi_mem, pi_bool, pi_map_flags, size_t, size_t, pi_uint32, const pi_event *, pi_event *, void **) { - // TODO : intel/llvm_test_suite DIE_NO_IMPLEMENTATION; } @@ -1625,7 +1620,6 @@ pi_result piextKernelSetArgPointer(pi_kernel, pi_uint32, size_t, const void *) { pi_result piextUSMEnqueueMemset(pi_queue, void *, pi_int32, size_t, pi_uint32, const pi_event *, pi_event *) { - // TODO : intel/llvm_test_suite DIE_NO_IMPLEMENTATION; } @@ -1636,7 +1630,6 @@ pi_result piextUSMEnqueueMemcpy(pi_queue, pi_bool, void *, const void *, size_t, pi_result piextUSMEnqueueMemAdvise(pi_queue, const void *, size_t, pi_mem_advice, pi_event *) { - // TODO : intel/llvm_test_suite DIE_NO_IMPLEMENTATION; } @@ -1674,8 +1667,7 @@ pi_result piextDeviceSelectBinary(pi_device, pi_device_binary *, pi_result piextUSMEnqueuePrefetch(pi_queue, const void *, size_t, pi_usm_migration_flags, pi_uint32, const pi_event *, pi_event *) { - // NOP for prefetch - return PI_SUCCESS; + DIE_NO_IMPLEMENTATION; } pi_result piextPluginGetOpaqueData(void *, void **OpaqueDataReturn) { From 11155d1514792b2b71efea2b03da67d51204606b Mon Sep 17 00:00:00 2001 From: Dongkyun Ahn Date: Fri, 17 Dec 2021 11:21:54 -0800 Subject: [PATCH 08/18] Build error fix after pulling sycl branch --- sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp | 7 +++++++ sycl/source/handler.cpp | 3 ++- 2 files changed, 9 insertions(+), 1 deletion(-) diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 4fc717cbaa6e8..4ae60ff52636f 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -865,6 +865,13 @@ pi_result piQueueFinish(pi_queue) { CONTINUE_NO_IMPLEMENTATION; } +pi_result piQueueFlush(pi_queue) { + // No-op as enqueued commands with ESIMD_EMULATOR plugin are blocking + // ones that do not return until their completion - kernel execution + // and memory read. + CONTINUE_NO_IMPLEMENTATION; +} + pi_result piextQueueGetNativeHandle(pi_queue, pi_native_handle *) { DIE_NO_IMPLEMENTATION; } diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index d0dabf790550a..7b45382c3161c 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -212,11 +212,12 @@ event handler::finalize() { backend::ext_intel_esimd_emulator) { // Dims==0 for 'single_task() - void(void) type' uint32_t Dims = (MArgs.size() > 0) ? MNDRDesc.Dims : 0; - return MQueue->getPlugin() + MQueue->getPlugin() .call( nullptr, reinterpret_cast(MHostKernel->getPtr()), Dims, &MNDRDesc.GlobalOffset[0], &MNDRDesc.GlobalSize[0], &MNDRDesc.LocalSize[0], 0, nullptr, nullptr); + return CL_SUCCESS; } else { return enqueueImpKernel(MQueue, MNDRDesc, MArgs, KernelBundleImpPtr, MKernel, MKernelName, MOSModuleHandle, From 7efcbc006748eb3a224f16edf3ff90da7af335ed Mon Sep 17 00:00:00 2001 From: Dongkyun Ahn Date: Fri, 17 Dec 2021 11:25:34 -0800 Subject: [PATCH 09/18] clang-format fix --- sycl/source/handler.cpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 7b45382c3161c..96c6b97ee6e01 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -212,11 +212,10 @@ event handler::finalize() { backend::ext_intel_esimd_emulator) { // Dims==0 for 'single_task() - void(void) type' uint32_t Dims = (MArgs.size() > 0) ? MNDRDesc.Dims : 0; - MQueue->getPlugin() - .call( - nullptr, reinterpret_cast(MHostKernel->getPtr()), - Dims, &MNDRDesc.GlobalOffset[0], &MNDRDesc.GlobalSize[0], - &MNDRDesc.LocalSize[0], 0, nullptr, nullptr); + MQueue->getPlugin().call( + nullptr, reinterpret_cast(MHostKernel->getPtr()), Dims, + &MNDRDesc.GlobalOffset[0], &MNDRDesc.GlobalSize[0], + &MNDRDesc.LocalSize[0], 0, nullptr, nullptr); return CL_SUCCESS; } else { return enqueueImpKernel(MQueue, MNDRDesc, MArgs, KernelBundleImpPtr, From 547176918dca40c4c2cd6da38db2e6c42b13951f Mon Sep 17 00:00:00 2001 From: Dongkyun Ahn Date: Mon, 20 Dec 2021 14:54:53 -0800 Subject: [PATCH 10/18] CM-EMU Device version info composition - Using version info fetched from calling 'CreateCmDevice' --- .../esimd_emulator/pi_esimd_emulator.cpp | 24 ++++++++++++++++--- 1 file changed, 21 insertions(+), 3 deletions(-) diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 4ae60ff52636f..325fa1a075b3f 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -127,6 +127,9 @@ static sycl::detail::ESIMDEmuPluginOpaqueData *PiESimdDeviceAccess; // For PI_DEVICE_INFO_DRIVER_VERSION info static char ESimdEmuVersionString[32]; +// For PI_DEVICE_INFO_VERSION info +static char CmEmuDeviceVersionString[32]; + using IDBuilder = sycl::detail::Builder; template @@ -463,6 +466,23 @@ pi_result piDevicesGet(pi_platform Platform, pi_device_type DeviceType, int Result = cm_support::CreateCmDevice(CmDevice, Version); + // CM Device version info consists of two decimal numbers - major + // and minor. Minor is single-digit. Version info is encoded into a + // unsigned integer value = 100 * major + minor. Second from right + // digit in decimal must be zero as it is used as 'dot' + // REF - $CM_EMU/common/cm_version_defs.h - 'CURRENT_CM_VERSION' + // e.g. CM version 7.3 => Device version = 703 + + if (((Version / 10) % 10) == 0) { + if (PrintPiTrace) { + std::cerr << "CM_EMU Device version info is incorrect" << std::endl; + } + return PI_INVALID_DEVICE; + } + + sprintf(CmEmuDeviceVersionString, "%d.%d", (int)(Version / 100), + (int)(Version % 10)); + if (Result != cm_support::CM_SUCCESS) { return PI_INVALID_DEVICE; } @@ -536,9 +556,7 @@ pi_result piDeviceGetInfo(pi_device Device, pi_device_info ParamName, // cl_khr_int64_extended_atomics return ReturnValue(""); case PI_DEVICE_INFO_VERSION: - // CM_EMU release version from - // https://github.com/intel/cm-cpu-emulation/releases - return ReturnValue("1.0"); + return ReturnValue(CmEmuDeviceVersionString); case PI_DEVICE_INFO_COMPILER_AVAILABLE: return ReturnValue(pi_bool{false}); case PI_DEVICE_INFO_LINKER_AVAILABLE: From 706c8f6461b7d332e99e3d19e2fda4adf7834d07 Mon Sep 17 00:00:00 2001 From: Dongkyun Ahn Date: Tue, 11 Jan 2022 17:15:08 -0800 Subject: [PATCH 11/18] Correction : Version format check --- sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 325fa1a075b3f..0b72ad21ea914 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -473,9 +473,9 @@ pi_result piDevicesGet(pi_platform Platform, pi_device_type DeviceType, // REF - $CM_EMU/common/cm_version_defs.h - 'CURRENT_CM_VERSION' // e.g. CM version 7.3 => Device version = 703 - if (((Version / 10) % 10) == 0) { + if (((Version / 10) % 10) != 0) { if (PrintPiTrace) { - std::cerr << "CM_EMU Device version info is incorrect" << std::endl; + std::cerr << "CM_EMU Device version info is incorrect : " << Version << std::endl; } return PI_INVALID_DEVICE; } From 02ebd7448c257696373e987e38878125ef4fe593 Mon Sep 17 00:00:00 2001 From: Dongkyun Ahn Date: Wed, 12 Jan 2022 09:17:48 -0800 Subject: [PATCH 12/18] clang-format error fix --- sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 0b72ad21ea914..08ec216925ec2 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -475,7 +475,8 @@ pi_result piDevicesGet(pi_platform Platform, pi_device_type DeviceType, if (((Version / 10) % 10) != 0) { if (PrintPiTrace) { - std::cerr << "CM_EMU Device version info is incorrect : " << Version << std::endl; + std::cerr << "CM_EMU Device version info is incorrect : " << Version + << std::endl; } return PI_INVALID_DEVICE; } From 9cd3b6deb07a4ee959aeaa2e6bd614db12b13912 Mon Sep 17 00:00:00 2001 From: Dongkyun Ahn Date: Thu, 13 Jan 2022 10:03:45 -0800 Subject: [PATCH 13/18] Patching works for suppressing warning message - Log messages from CM_EMU causes failures from intel/llvm-test-suite - Some CM_EMU's files are patched before library building to suppress such log messages with "PATCH_COMMAND" in CMake file --- sycl/plugins/esimd_emulator/CMakeLists.txt | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/sycl/plugins/esimd_emulator/CMakeLists.txt b/sycl/plugins/esimd_emulator/CMakeLists.txt index c3660bc2bcd0d..7c11c7464445e 100755 --- a/sycl/plugins/esimd_emulator/CMakeLists.txt +++ b/sycl/plugins/esimd_emulator/CMakeLists.txt @@ -63,10 +63,18 @@ else () if (MSVC) message(FATAL_ERROR "Online-building of CM_EMU library is not supported under Windows environment") else() + # Arguments for online patching to suppress log message from CM_EMU + # 1. Replacing CM_EMU's log print-out macro controlled by 'GFX_EMU_WITH_FLAGS_' + # with blank space from $CM_EMU_SRC/common/emu_log.h + set (replacing_pattern s/{\ ?GFX_EMU_WITH_FLAGS_.*//g) + # 2. Range of lines to be removed printing out unnecessary log messages + # from $CM_EMU_SRC/libcmrt/cm_device_emumode.cpp + set (removed_range 130,132d) ExternalProject_Add(cm-emu GIT_REPOSITORY https://github.com/intel/cm-cpu-emulation.git BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}/cm-emu_build INSTALL_DIR ${CMAKE_CURRENT_BINARY_DIR}/cm-emu_install + PATCH_COMMAND perl -pi.back -e ${replacing_pattern} ${CMAKE_CURRENT_BINARY_DIR}/cm-emu-prefix/src/cm-emu/common/emu_log.h && sed --in-place ${removed_range} ${CMAKE_CURRENT_BINARY_DIR}/cm-emu-prefix/src/cm-emu/libcmrt/cm_device_emumode.cpp CMAKE_ARGS -DLIBVA_INSTALL_PATH=/usr -D__SYCL_EXPLICIT_SIMD_PLUGIN__=true -DCMAKE_INSTALL_PREFIX= From d0ac755a213793dd3d56e9fa683a3687e31ced42 Mon Sep 17 00:00:00 2001 From: Dongkyun Ahn Date: Fri, 14 Jan 2022 12:08:28 -0800 Subject: [PATCH 14/18] Returning single-entry CM device list only for GPU type request --- .../esimd_emulator/pi_esimd_emulator.cpp | 20 ++++++++++++++++++- 1 file changed, 19 insertions(+), 1 deletion(-) diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 08ec216925ec2..f52e9a08ed2c4 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -449,8 +449,26 @@ pi_result piDevicesGet(pi_platform Platform, pi_device_type DeviceType, } // CM has single-root-device without sub-device support. + pi_uint32 DeviceCount = + (DeviceType == PI_DEVICE_TYPE_GPU || DeviceType == PI_DEVICE_TYPE_DEFAULT) + ? 1 + : 0; + if (NumDevices) { - *NumDevices = 1; + *NumDevices = DeviceCount; + } + + if (NumEntries == 0) { + /// Runtime queries number of devices + assert(Devices == nullptr && + "Devices should be nullptr when querying the number of devices"); + + return PI_SUCCESS; + } + + if (DeviceCount == 0) { + /// No GPU entry to fill 'Device' array + return PI_SUCCESS; } cm_support::CmDevice *CmDevice = nullptr; From d71a04d56cacc3fc7fbb839f74313a86a974a51d Mon Sep 17 00:00:00 2001 From: Dongkyun Ahn Date: Fri, 14 Jan 2022 17:01:28 -0800 Subject: [PATCH 15/18] Updating returning platform info with NumEntries And, - Argument sanity check failure revised for piDevicesGet - DeviceType condition check fix --- .../esimd_emulator/pi_esimd_emulator.cpp | 40 ++++++++++++------- 1 file changed, 26 insertions(+), 14 deletions(-) diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index f52e9a08ed2c4..7fa02763602d7 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -382,9 +382,23 @@ pi_result piPlatformsGet(pi_uint32 NumEntries, pi_platform *Platforms, PrintPiTrace = true; } - if (NumEntries == 0 && Platforms != nullptr) { - return PI_INVALID_VALUE; + if (NumPlatforms) { + *NumPlatforms = 1; + } + + if (NumEntries == 0) { + /// Runtime queries number of Platforms + if (Platforms != nullptr) { + if (PrintPiTrace) { + std::cerr << "Invalid Arguments for piPlatformsGet of esimd_emultor " + "(Platforms!=nullptr) while querying number of platforms" + << std::endl; + } + return PI_INVALID_VALUE; + } + return PI_SUCCESS; } + if (Platforms == nullptr && NumPlatforms == nullptr) { return PI_INVALID_VALUE; } @@ -394,10 +408,6 @@ pi_result piPlatformsGet(pi_uint32 NumEntries, pi_platform *Platforms, Platforms[0]->CmEmuVersion = std::string("0.0.1"); } - if (NumPlatforms) { - *NumPlatforms = 1; - } - return PI_SUCCESS; } @@ -448,11 +458,8 @@ pi_result piDevicesGet(pi_platform Platform, pi_device_type DeviceType, return PI_INVALID_PLATFORM; } - // CM has single-root-device without sub-device support. - pi_uint32 DeviceCount = - (DeviceType == PI_DEVICE_TYPE_GPU || DeviceType == PI_DEVICE_TYPE_DEFAULT) - ? 1 - : 0; + // CM has single-root-GPU-device without sub-device support. + pi_uint32 DeviceCount = (DeviceType & PI_DEVICE_TYPE_GPU) ? 1 : 0; if (NumDevices) { *NumDevices = DeviceCount; @@ -460,9 +467,14 @@ pi_result piDevicesGet(pi_platform Platform, pi_device_type DeviceType, if (NumEntries == 0) { /// Runtime queries number of devices - assert(Devices == nullptr && - "Devices should be nullptr when querying the number of devices"); - + if (Devices != nullptr) { + if (PrintPiTrace) { + std::cerr << "Invalid Arguments for piDevicesGet of esimd_emultor " + "(Devices!=nullptr) while querying number of platforms" + << std::endl; + } + return PI_INVALID_VALUE; + } return PI_SUCCESS; } From 6f1034e9c2e1bcb6f729791e7b2eaa17538851d4 Mon Sep 17 00:00:00 2001 From: Dongkyun Ahn Date: Sat, 15 Jan 2022 10:49:55 -0800 Subject: [PATCH 16/18] Single-point-of-return for EnqueueKernel --- sycl/source/handler.cpp | 14 +++++++++----- 1 file changed, 9 insertions(+), 5 deletions(-) diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index dc9767b3bb401..3cdf937ba4dfc 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -203,10 +203,12 @@ event handler::finalize() { RT::PiEvent *OutEvent = nullptr; auto EnqueueKernel = [&]() { + // 'Result' for single point of return + cl_int Result = CL_INVALID_VALUE; if (MQueue->is_host()) { MHostKernel->call( MNDRDesc, (NewEvent) ? NewEvent->getHostProfilingInfo() : nullptr); - return CL_SUCCESS; + Result = CL_SUCCESS; } else { if (MQueue->getPlugin().getBackend() == backend::ext_intel_esimd_emulator) { @@ -216,13 +218,15 @@ event handler::finalize() { nullptr, reinterpret_cast(MHostKernel->getPtr()), Dims, &MNDRDesc.GlobalOffset[0], &MNDRDesc.GlobalSize[0], &MNDRDesc.LocalSize[0], 0, nullptr, nullptr); - return CL_SUCCESS; + Result = CL_SUCCESS; } else { - return enqueueImpKernel(MQueue, MNDRDesc, MArgs, KernelBundleImpPtr, - MKernel, MKernelName, MOSModuleHandle, - RawEvents, OutEvent, nullptr); + Result = enqueueImpKernel(MQueue, MNDRDesc, MArgs, KernelBundleImpPtr, + MKernel, MKernelName, MOSModuleHandle, + RawEvents, OutEvent, nullptr); } } + // assert(Result != CL_INVALID_VALUE); + return Result; }; bool DiscardEvent = false; From 71802cc8220436294810e2db73a75833d5754c3a Mon Sep 17 00:00:00 2001 From: Dongkyun Ahn Date: Tue, 18 Jan 2022 12:23:13 -0800 Subject: [PATCH 17/18] Adding command line option for disabling ESIMD_EMULATOR build --- buildbot/configure.py | 6 ++++++ sycl/plugins/CMakeLists.txt | 10 ++++++---- 2 files changed, 12 insertions(+), 4 deletions(-) diff --git a/buildbot/configure.py b/buildbot/configure.py index 3ee40e2cd6188..4be8509409144 100644 --- a/buildbot/configure.py +++ b/buildbot/configure.py @@ -32,6 +32,7 @@ def do_configure(args): libclc_targets_to_build = '' libclc_gen_remangled_variants = 'OFF' sycl_build_pi_cuda = 'OFF' + sycl_build_pi_esimd_emulator = 'ON' sycl_build_pi_hip = 'OFF' sycl_build_pi_hip_platform = 'AMD' sycl_clang_extra_flags = '' @@ -49,6 +50,9 @@ def do_configure(args): if args.arm: llvm_targets_to_build = 'ARM;AArch64' + if args.disable_esimd_emulator: + sycl_build_pi_esimd_emulator = 'OFF' + if args.cuda or args.hip: llvm_enable_projects += ';libclc' @@ -143,6 +147,7 @@ def do_configure(args): "-DBUILD_SHARED_LIBS={}".format(llvm_build_shared_libs), "-DSYCL_ENABLE_XPTI_TRACING={}".format(sycl_enable_xpti_tracing), "-DLLVM_ENABLE_LLD={}".format(llvm_enable_lld), + "-DSYCL_BUILD_PI_ESIMD_EMULATOR={}".format(sycl_build_pi_esimd_emulator), "-DXPTI_ENABLE_WERROR={}".format(xpti_enable_werror), "-DSYCL_CLANG_EXTRA_FLAGS={}".format(sycl_clang_extra_flags) ] @@ -208,6 +213,7 @@ def main(): parser.add_argument("--hip-platform", type=str, choices=['AMD', 'NVIDIA'], default='AMD', help="choose hardware platform for HIP backend") parser.add_argument("--hip-amd-arch", type=str, help="Sets AMD gpu architecture for llvm lit tests, this is only needed for the HIP backend and AMD platform") parser.add_argument("--arm", action='store_true', help="build ARM support rather than x86") + parser.add_argument("--disable-esimd-emulator", action='store_true', help="exclude ESIMD_EMULATOR 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") diff --git a/sycl/plugins/CMakeLists.txt b/sycl/plugins/CMakeLists.txt index 25ce821b138c6..12fd21881f158 100644 --- a/sycl/plugins/CMakeLists.txt +++ b/sycl/plugins/CMakeLists.txt @@ -15,8 +15,10 @@ endif() add_subdirectory(opencl) add_subdirectory(level_zero) -# TODO : Remove 'if (NOT MSVC)' when CM_EMU supports Windows -# environment -if (NOT MSVC) - add_subdirectory(esimd_emulator) +if(SYCL_BUILD_PI_ESIMD_EMULATOR) + # TODO : Remove 'if (NOT MSVC)' when CM_EMU supports Windows + # environment + if (NOT MSVC) + add_subdirectory(esimd_emulator) + endif() endif() From 139a54975f35c0b3a7597b630b5470339271f197 Mon Sep 17 00:00:00 2001 From: Dongkyun Ahn Date: Tue, 18 Jan 2022 12:57:09 -0800 Subject: [PATCH 18/18] Updating pi_esimd_emulator build due to update in open-source CM_EMU - Log-in suppression - 'cpu' to 'emu' (PR #4728) --- sycl/plugins/esimd_emulator/CMakeLists.txt | 10 ++++------ sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp | 6 ++---- 2 files changed, 6 insertions(+), 10 deletions(-) diff --git a/sycl/plugins/esimd_emulator/CMakeLists.txt b/sycl/plugins/esimd_emulator/CMakeLists.txt index 7c11c7464445e..fc0196193d74e 100755 --- a/sycl/plugins/esimd_emulator/CMakeLists.txt +++ b/sycl/plugins/esimd_emulator/CMakeLists.txt @@ -64,17 +64,15 @@ else () message(FATAL_ERROR "Online-building of CM_EMU library is not supported under Windows environment") else() # Arguments for online patching to suppress log message from CM_EMU - # 1. Replacing CM_EMU's log print-out macro controlled by 'GFX_EMU_WITH_FLAGS_' - # with blank space from $CM_EMU_SRC/common/emu_log.h + # Replacing CM_EMU's log print-out macro controlled by 'GFX_EMU_WITH_FLAGS_' + # with blank space from $CM_EMU_SRC/common/emu_log.h set (replacing_pattern s/{\ ?GFX_EMU_WITH_FLAGS_.*//g) - # 2. Range of lines to be removed printing out unnecessary log messages - # from $CM_EMU_SRC/libcmrt/cm_device_emumode.cpp - set (removed_range 130,132d) ExternalProject_Add(cm-emu GIT_REPOSITORY https://github.com/intel/cm-cpu-emulation.git + GIT_TAG c19234cea13bdfc32b5ed9 BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}/cm-emu_build INSTALL_DIR ${CMAKE_CURRENT_BINARY_DIR}/cm-emu_install - PATCH_COMMAND perl -pi.back -e ${replacing_pattern} ${CMAKE_CURRENT_BINARY_DIR}/cm-emu-prefix/src/cm-emu/common/emu_log.h && sed --in-place ${removed_range} ${CMAKE_CURRENT_BINARY_DIR}/cm-emu-prefix/src/cm-emu/libcmrt/cm_device_emumode.cpp + PATCH_COMMAND perl -pi.back -e ${replacing_pattern} ${CMAKE_CURRENT_BINARY_DIR}/cm-emu-prefix/src/cm-emu/common/emu_log.h CMAKE_ARGS -DLIBVA_INSTALL_PATH=/usr -D__SYCL_EXPLICIT_SIMD_PLUGIN__=true -DCMAKE_INSTALL_PREFIX= diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 7fa02763602d7..363d5e86c14fa 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -29,7 +29,7 @@ #include // TODO : Rename esimdcpu to esimdemu for next CM_EMU release -#include +#include #include #include @@ -249,9 +249,7 @@ template class libCMBatch { GroupDim[I] = (uint32_t)(GlobalSize[I] / LocalSize[I]); } - // TODO : Replace ESimdCPUKernel with EsimdemuKernel for next - // CM_EMU release - ESimdCPUKernel Esimdemu((fptrVoid)InvokeLambda, GroupDim, SpaceDim); + EsimdemuKernel Esimdemu((fptrVoid)InvokeLambda, GroupDim, SpaceDim); Esimdemu.launchMT(sizeof(struct LambdaWrapper), WrappedLambda.get()); }