diff --git a/sycl/CMakeLists.txt b/sycl/CMakeLists.txt index 87883dc217969..8892251639f1a 100644 --- a/sycl/CMakeLists.txt +++ b/sycl/CMakeLists.txt @@ -321,17 +321,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. diff --git a/sycl/include/CL/sycl/detail/cg_types.hpp b/sycl/include/CL/sycl/detail/cg_types.hpp index 24ec70f9e9a9e..fe5e52f0d2ab9 100644 --- a/sycl/include/CL/sycl/detail/cg_types.hpp +++ b/sycl/include/CL/sycl/detail/cg_types.hpp @@ -246,6 +246,9 @@ template 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) {} diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index ab0925e197e78..c3ddbc1930258 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/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 + KernelType *ResetHostKernelHelper(const KernelType &KernelFunc) { + NormalizedKernelType NormalizedKernel(KernelFunc); + auto NormalizedKernelFunc = + std::function &)>(NormalizedKernel); + auto HostKernelPtr = + new detail::HostKernel, Dims, KernelName>( + NormalizedKernelFunc); + MHostKernel.reset(HostKernelPtr); + return &HostKernelPtr->MKernel.template target() + ->MKernelFunc; + } + + // For 'sycl::id' kernel argument + template + typename std::enable_if>::value, + KernelType *>::type + ResetHostKernel(const KernelType &KernelFunc) { + struct NormalizedKernelType { + KernelType MKernelFunc; + NormalizedKernelType(const KernelType &KernelFunc) + : MKernelFunc(KernelFunc) {} + void operator()(const nd_item &Arg) { + detail::runKernelWithArg(MKernelFunc, Arg.get_global_id()); + } + }; + return ResetHostKernelHelper(KernelFunc); + } + + // For 'sycl::nd_item' kernel argument + template + typename std::enable_if>::value, + KernelType *>::type + ResetHostKernel(const KernelType &KernelFunc) { + struct NormalizedKernelType { + KernelType MKernelFunc; + NormalizedKernelType(const KernelType &KernelFunc) + : MKernelFunc(KernelFunc) {} + void operator()(const nd_item &Arg) { + detail::runKernelWithArg(MKernelFunc, Arg); + } + }; + return ResetHostKernelHelper(KernelFunc); + } + + // For 'sycl::item' kernel argument + template + typename std::enable_if>::value, + KernelType *>::type + ResetHostKernel(const KernelType &KernelFunc) { + struct NormalizedKernelType { + KernelType MKernelFunc; + NormalizedKernelType(const KernelType &KernelFunc) + : MKernelFunc(KernelFunc) {} + void operator()(const nd_item &Arg) { + sycl::item Item = detail::Builder::createItem( + Arg.get_global_range(), Arg.get_global_id()); + detail::runKernelWithArg(MKernelFunc, Item); + } + }; + return ResetHostKernelHelper(KernelFunc); + } + + // For 'sycl::item' kernel argument + template + typename std::enable_if>::value, + KernelType *>::type + ResetHostKernel(const KernelType &KernelFunc) { + struct NormalizedKernelType { + KernelType MKernelFunc; + NormalizedKernelType(const KernelType &KernelFunc) + : MKernelFunc(KernelFunc) {} + void operator()(const nd_item &Arg) { + sycl::item Item = detail::Builder::createItem( + Arg.get_global_range(), Arg.get_global_id(), Arg.get_offset()); + detail::runKernelWithArg(MKernelFunc, Item); + } + }; + return ResetHostKernelHelper(KernelFunc); + } + + /* 'wrapper'-based approach using 'NormalizedKernelType' struct is + * not applied for 'void(void)' type kernel and + * 'void(sycl::group)'. This is because 'void(void)' type does + * not have argument to normalize and 'void(sycl::group)' is + * not supported in ESIMD. + */ + // For 'void' and 'sycl::group' kernel argument + template + typename std::enable_if::value || + std::is_same>::value, + KernelType *>::type + ResetHostKernel(const KernelType &KernelFunc) { + MHostKernel.reset( + new detail::HostKernel(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,9 +706,9 @@ class __SYCL_EXPORT handler { "kernel_handler is not yet supported by host device.", PI_INVALID_OPERATION); } - MHostKernel.reset( - new detail::HostKernel( - KernelFunc)); + KernelType *KernelPtr = + ResetHostKernel( + KernelFunc); using KI = sycl::detail::KernelInfo; // Empty name indicates that the compilation happens without integration @@ -598,8 +716,9 @@ class __SYCL_EXPORT handler { 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(KernelPtr), + KI::getNumParams(), &KI::getParamDesc(0), + KI::isESIMD()); MKernelName = KI::getName(); MOSModuleHandle = detail::OSUtil::getOSModuleHandle(KI::getName()); } else { diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/detail/atomic_intrin.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/detail/atomic_intrin.hpp new file mode 100644 index 0000000000000..d52074137ef20 --- /dev/null +++ b/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 + +// This function implements atomic update of pre-existing variable in the +// absense of C++ 20's atomic_ref. +template 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 +} diff --git a/sycl/include/sycl/ext/intel/experimental/esimd/emu/detail/esimd_emulator_device_interface.hpp b/sycl/include/sycl/ext/intel/experimental/esimd/emu/detail/esimd_emulator_device_interface.hpp index 4a57576c1753b..5f539e89ee67f 100644 --- a/sycl/include/sycl/ext/intel/experimental/esimd/emu/detail/esimd_emulator_device_interface.hpp +++ b/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(nullptr); - - ESIMDEmuPluginOpaqueData *OpaqueData = - reinterpret_cast(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(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 diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index ddaef3273f831..5f614f57d43bb 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -176,6 +176,7 @@ set(SYCL_SOURCES "sampler.cpp" "stream.cpp" "spirv_ops.cpp" + "esimd_emulator_device_interface.cpp" "$<$:detail/windows_pi.cpp>" "$<$,$>:detail/posix_pi.cpp>" "$<$:abi_replacements_windows.cpp>" diff --git a/sycl/source/detail/config.cpp b/sycl/source/detail/config.cpp index 7c02439b2d9fb..4ac9700a5e48c 100644 --- a/sycl/source/detail/config.cpp +++ b/sycl/source/detail/config.cpp @@ -169,13 +169,14 @@ getSyclDeviceTypeMap() { } // Array is used by SYCL_DEVICE_FILTER and SYCL_DEVICE_ALLOWLIST -const std::array, 6> &getSyclBeMap() { - static const std::array, 6> SyclBeMap = { +const std::array, 7> &getSyclBeMap() { + static const std::array, 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; } diff --git a/sycl/source/detail/config.hpp b/sycl/source/detail/config.hpp index 8589d5b96899f..fe9f1de0ce3f5 100644 --- a/sycl/source/detail/config.hpp +++ b/sycl/source/detail/config.hpp @@ -134,12 +134,13 @@ template <> class SYCLConfig { return BackendPtr; const char *ValStr = BaseT::getRawValue(); - const std::array, 5> SyclBeMap = { + const std::array, 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 { }); 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, 5> & getSyclDeviceTypeMap(); // Array is used by SYCL_DEVICE_FILTER and SYCL_DEVICE_ALLOWLIST -const std::array, 6> &getSyclBeMap(); +const std::array, 7> &getSyclBeMap(); template <> class SYCLConfig { using BaseT = SYCLConfigBase; diff --git a/sycl/source/detail/device_filter.cpp b/sycl/source/detail/device_filter.cpp index 6017996e8d40b..6b3746e7c57ac 100644 --- a/sycl/source/detail/device_filter.cpp +++ b/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); diff --git a/sycl/source/detail/pi.cpp b/sycl/source/detail/pi.cpp index f7127e1421e28..316ab7d432a98 100644 --- a/sycl/source/detail/pi.cpp +++ b/sycl/source/detail/pi.cpp @@ -285,11 +285,14 @@ std::vector> 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 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> 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 &Plugins) { // Use the LEVEL_ZERO plugin as the GlobalPlugin GlobalPlugin = std::make_shared( 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( + PluginInformation, backend::ext_intel_esimd_emulator, Library); } Plugins.emplace_back( plugin(PluginInformation, PluginNames[I].second, Library)); diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index 65dfcc285f75b..b38c1894daa40 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -2134,7 +2134,8 @@ cl_int ExecCGCommand::enqueueImp() { NDRDescT &NDRDesc = ExecKernel->MNDRDesc; std::vector &Args = ExecKernel->MArgs; - if (MQueue->is_host()) { + if (MQueue->is_host() || (MQueue->getPlugin().getBackend() == + backend::ext_intel_esimd_emulator)) { for (ArgDesc &Arg : Args) if (kernel_param_kind_t::kind_accessor == Arg.MType) { Requirement *Req = (Requirement *)(Arg.MPtr); @@ -2146,8 +2147,19 @@ cl_int ExecCGCommand::enqueueImp() { const detail::plugin &Plugin = EventImpls[0]->getPlugin(); Plugin.call(RawEvents.size(), &RawEvents[0]); } - ExecKernel->MHostKernel->call(NDRDesc, - getEvent()->getHostProfilingInfo()); + + if (MQueue->is_host()) { + ExecKernel->MHostKernel->call(NDRDesc, + getEvent()->getHostProfilingInfo()); + } else { + assert(MQueue->getPlugin().getBackend() == + backend::ext_intel_esimd_emulator); + MQueue->getPlugin().call( + nullptr, + reinterpret_cast(ExecKernel->MHostKernel->getPtr()), + NDRDesc.Dims, &NDRDesc.GlobalOffset[0], &NDRDesc.GlobalSize[0], + &NDRDesc.LocalSize[0], 0, nullptr, nullptr); + } return CL_SUCCESS; } diff --git a/sycl/source/esimd_emulator_device_interface.cpp b/sycl/source/esimd_emulator_device_interface.cpp new file mode 100644 index 0000000000000..ad9f54b8b5a75 --- /dev/null +++ b/sycl/source/esimd_emulator_device_interface.cpp @@ -0,0 +1,70 @@ +//==--------------- esimd_emulator_device_interface.cpp --------------------==// +// +// 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.cpp +/// Definitions for ESIMD_EMULATOR-device specific definitions. +/// +/// This interface is for ESIMD intrinsic emulation implementations +/// such as slm_access to access ESIMD_EMULATOR specific-support therefore +/// it has to be defined and shared as include directory +/// +/// \ingroup sycl_pi_esimd_emulator + +#include +#include + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace detail { + +__SYCL_EXPORT 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(nullptr); + + ESIMDEmuPluginOpaqueData *OpaqueData = + reinterpret_cast(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 feature_not_supported(); + } + // Opaque data version is OK, can cast the 'data' field. + ESIMDDeviceInterface *Interface = + reinterpret_cast(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 feature_not_supported(); + } + return Interface; +} +} // namespace detail +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 95e0d29106179..7d0487da9eb9f 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3854,6 +3854,7 @@ _ZN2cl4sycl6detail22get_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6devic _ZN2cl4sycl6detail22get_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6deviceESaIS6_EERKS5_INS0_9kernel_idESaISB_EENS0_12bundle_stateE _ZN2cl4sycl6detail22has_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6deviceESaIS6_EENS0_12bundle_stateE _ZN2cl4sycl6detail22has_kernel_bundle_implERKNS0_7contextERKSt6vectorINS0_6deviceESaIS6_EERKS5_INS0_9kernel_idESaISB_EENS0_12bundle_stateE +_ZN2cl4sycl6detail23getESIMDDeviceInterfaceEv _ZN2cl4sycl6detail24find_device_intersectionERKSt6vectorINS0_13kernel_bundleILNS0_12bundle_stateE1EEESaIS5_EE _ZN2cl4sycl6detail27getPixelCoordLinearFiltModeENS0_3vecIfLi4EEENS0_15addressing_modeENS0_5rangeILi3EEERS3_ _ZN2cl4sycl6detail28getDeviceFunctionPointerImplERNS0_6deviceERNS0_7programEPKc diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 0df42b5fa4e8d..0f2cebda7385a 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -1847,6 +1847,7 @@ ?fill@MemoryManager@detail@sycl@cl@@SAXPEAVSYCLMemObjI@234@PEAXV?$shared_ptr@Vqueue_impl@detail@sycl@cl@@@std@@_KPEBDIV?$range@$02@34@5V?$id@$02@34@IV?$vector@PEAU_pi_event@@V?$allocator@PEAU_pi_event@@@std@@@7@AEAPEAU_pi_event@@@Z ?fill_usm@MemoryManager@detail@sycl@cl@@SAXPEAXV?$shared_ptr@Vqueue_impl@detail@sycl@cl@@@std@@_KHV?$vector@PEAU_pi_event@@V?$allocator@PEAU_pi_event@@@std@@@6@AEAPEAU_pi_event@@@Z ?finalize@handler@sycl@cl@@AEAA?AVevent@23@XZ +?getESIMDDeviceInterface@detail@sycl@cl@@YAPEAUESIMDDeviceInterface@123@XZ ?find_device_intersection@detail@sycl@cl@@YA?AV?$vector@Vdevice@sycl@cl@@V?$allocator@Vdevice@sycl@cl@@@std@@@std@@AEBV?$vector@V?$kernel_bundle@$00@sycl@cl@@V?$allocator@V?$kernel_bundle@$00@sycl@cl@@@std@@@5@@Z ?floor@__host_std@cl@@YA?AV?$vec@M$00@sycl@2@V342@@Z ?floor@__host_std@cl@@YA?AV?$vec@M$01@sycl@2@V342@@Z diff --git a/sycl/unittests/allowlist/ParseAllowList.cpp b/sycl/unittests/allowlist/ParseAllowList.cpp index 637defd908088..0d2a8edb91ba7 100644 --- a/sycl/unittests/allowlist/ParseAllowList.cpp +++ b/sycl/unittests/allowlist/ParseAllowList.cpp @@ -167,7 +167,8 @@ TEST(ParseAllowListTests, CheckAllValidBackendNameValuesAreProcessed) { sycl::detail::AllowListParsedT ExpectedValue{ {{"BackendName", "host"}}, {{"BackendName", "opencl"}}, {{"BackendName", "level_zero"}}, {{"BackendName", "cuda"}}, - {{"BackendName", "hip"}}, {{"BackendName", "*"}}}; + {{"BackendName", "hip"}}, {{"BackendName", "esimd_emulator"}}, + {{"BackendName", "*"}}}; EXPECT_EQ(ExpectedValue, ActualValue); }