diff --git a/libdevice/atomic.hpp b/libdevice/atomic.hpp new file mode 100644 index 0000000000000..bc1188a3844f9 --- /dev/null +++ b/libdevice/atomic.hpp @@ -0,0 +1,95 @@ +//==-------------- atomic.hpp - support of atomic operations ---------------==// +// +// 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 + +#include "device.h" + +#ifdef __SPIR__ + +#define __SYCL_GLOBAL__ __attribute__((opencl_global)) + +namespace __spv { +struct Scope { + + enum Flag : uint32_t { + CrossDevice = 0, + Device = 1, + Workgroup = 2, + Subgroup = 3, + Invocation = 4, + }; + + constexpr Scope(Flag flag) : flag_value(flag) {} + + constexpr operator uint32_t() const { return flag_value; } + + Flag flag_value; +}; + +struct MemorySemanticsMask { + + enum Flag : uint32_t { + None = 0x0, + Acquire = 0x2, + Release = 0x4, + AcquireRelease = 0x8, + SequentiallyConsistent = 0x10, + UniformMemory = 0x40, + SubgroupMemory = 0x80, + WorkgroupMemory = 0x100, + CrossWorkgroupMemory = 0x200, + AtomicCounterMemory = 0x400, + ImageMemory = 0x800, + }; + + constexpr MemorySemanticsMask(Flag flag) : flag_value(flag) {} + + constexpr operator uint32_t() const { return flag_value; } + + Flag flag_value; +}; +} // namespace __spv + +extern DEVICE_EXTERNAL int +__spirv_AtomicCompareExchange(int __SYCL_GLOBAL__ *, __spv::Scope::Flag, + __spv::MemorySemanticsMask::Flag, + __spv::MemorySemanticsMask::Flag, int, int); + +extern DEVICE_EXTERNAL int __spirv_AtomicLoad(const int __SYCL_GLOBAL__ *, + __spv::Scope::Flag, + __spv::MemorySemanticsMask::Flag); + +extern DEVICE_EXTERNAL int __spirv_AtomicStore(int __SYCL_GLOBAL__ *, + __spv::Scope::Flag, + __spv::MemorySemanticsMask::Flag, + int); + +/// Atomically set the value in *Ptr with Desired if and only if it is Expected +/// Return the value which already was in *Ptr +static inline int atomicCompareAndSet(__SYCL_GLOBAL__ int *Ptr, int Desired, + int Expected) { + return __spirv_AtomicCompareExchange( + Ptr, __spv::Scope::Device, + __spv::MemorySemanticsMask::SequentiallyConsistent, + __spv::MemorySemanticsMask::SequentiallyConsistent, Desired, Expected); +} + +static inline int atomicLoad(__SYCL_GLOBAL__ int *Ptr) { + return __spirv_AtomicLoad(Ptr, __spv::Scope::Device, + __spv::MemorySemanticsMask::SequentiallyConsistent); +} + +static inline int atomicStore(__SYCL_GLOBAL__ int *Ptr, int V) { + return __spirv_AtomicStore(Ptr, __spv::Scope::Device, + __spv::MemorySemanticsMask::SequentiallyConsistent, + V); +} + +#endif // __SPIR__ diff --git a/libdevice/fallback-cassert.cpp b/libdevice/fallback-cassert.cpp index 724d4635fb0b5..fc42ac64c748a 100644 --- a/libdevice/fallback-cassert.cpp +++ b/libdevice/fallback-cassert.cpp @@ -6,26 +6,97 @@ // //===----------------------------------------------------------------------===// +#include "atomic.hpp" +#include "include/assert-happened.hpp" #include "wrapper.h" #ifdef __SPIR__ + +#define ASSERT_NONE 0 +#define ASSERT_START 1 +#define ASSERT_FINISH 2 + +// definition +__SYCL_GLOBAL__ AssertHappened __SYCL_AssertHappenedMem; + static const __attribute__((opencl_constant)) char assert_fmt[] = "%s:%d: %s: global id: [%lu,%lu,%lu], local id: [%lu,%lu,%lu] " "Assertion `%s` failed.\n"; +DEVICE_EXTERN_C void __devicelib_assert_read(void *_Dst) { + AssertHappened *Dst = (AssertHappened *)_Dst; + int Flag = atomicLoad(&__SYCL_AssertHappenedMem.Flag); + + if (ASSERT_NONE == Flag) { + Dst->Flag = Flag; + return; + } + + if (Flag != ASSERT_FINISH) + while (ASSERT_START == atomicLoad(&__SYCL_AssertHappenedMem.Flag)) + ; + + *Dst = __SYCL_AssertHappenedMem; +} + DEVICE_EXTERN_C void __devicelib_assert_fail(const char *expr, const char *file, int32_t line, const char *func, uint64_t gid0, uint64_t gid1, uint64_t gid2, uint64_t lid0, uint64_t lid1, uint64_t lid2) { - // intX_t types are used instead of `int' and `long' because the format string - // is defined in terms of *device* types (OpenCL types): %d matches a 32 bit - // integer, %lu matches a 64 bit unsigned integer. Host `int' and - // `long' types may be different, so we cannot use them. - __spirv_ocl_printf(assert_fmt, file, (int32_t)line, - // WORKAROUND: IGC does not handle this well - // (func) ? func : "", - func, gid0, gid1, gid2, lid0, lid1, lid2, expr); + int Expected = ASSERT_NONE; + int Desired = ASSERT_START; + + if (atomicCompareAndSet(&__SYCL_AssertHappenedMem.Flag, Desired, Expected) == + Expected) { + __SYCL_AssertHappenedMem.Line = line; + __SYCL_AssertHappenedMem.GID0 = gid0; + __SYCL_AssertHappenedMem.GID1 = gid1; + __SYCL_AssertHappenedMem.GID2 = gid2; + __SYCL_AssertHappenedMem.LID0 = lid0; + __SYCL_AssertHappenedMem.LID1 = lid1; + __SYCL_AssertHappenedMem.LID2 = lid2; + + int ExprLength = 0; + int FileLength = 0; + int FuncLength = 0; + + if (expr) + for (const char *C = expr; *C != '\0'; ++C, ++ExprLength) + ; + if (file) + for (const char *C = file; *C != '\0'; ++C, ++FileLength) + ; + if (func) + for (const char *C = func; *C != '\0'; ++C, ++FuncLength) + ; + + int MaxExprIdx = sizeof(__SYCL_AssertHappenedMem.Expr) - 1; + int MaxFileIdx = sizeof(__SYCL_AssertHappenedMem.File) - 1; + int MaxFuncIdx = sizeof(__SYCL_AssertHappenedMem.Func) - 1; + + if (ExprLength < MaxExprIdx) + MaxExprIdx = ExprLength; + if (FileLength < MaxFileIdx) + MaxFileIdx = FileLength; + if (FuncLength < MaxFuncIdx) + MaxFuncIdx = FuncLength; + + for (int Idx = 0; Idx < MaxExprIdx; ++Idx) + __SYCL_AssertHappenedMem.Expr[Idx] = expr[Idx]; + __SYCL_AssertHappenedMem.Expr[MaxExprIdx] = '\0'; + + for (int Idx = 0; Idx < MaxFileIdx; ++Idx) + __SYCL_AssertHappenedMem.File[Idx] = file[Idx]; + __SYCL_AssertHappenedMem.File[MaxFileIdx] = '\0'; + + for (int Idx = 0; Idx < MaxFuncIdx; ++Idx) + __SYCL_AssertHappenedMem.Func[Idx] = func[Idx]; + __SYCL_AssertHappenedMem.Func[MaxFuncIdx] = '\0'; + + // Show we've done copying + atomicStore(&__SYCL_AssertHappenedMem.Flag, ASSERT_FINISH); + } // FIXME: call SPIR-V unreachable instead // volatile int *die = (int *)0x0; diff --git a/libdevice/include/assert-happened.hpp b/libdevice/include/assert-happened.hpp new file mode 100644 index 0000000000000..a4d53f6401240 --- /dev/null +++ b/libdevice/include/assert-happened.hpp @@ -0,0 +1,42 @@ +//==-- assert-happened.hpp - Structure and declaration for assert support --==// +// +// 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 + +// Treat this header as system one to workaround frontend's restriction +#pragma clang system_header + +#ifdef __SPIR__ + +struct AssertHappened { + int Flag = 0; + char Expr[256 + 1] = ""; + char File[256 + 1] = ""; + char Func[128 + 1] = ""; + + int32_t Line = 0; + + uint64_t GID0 = 0; + uint64_t GID1 = 0; + uint64_t GID2 = 0; + + uint64_t LID0 = 0; + uint64_t LID1 = 0; + uint64_t LID2 = 0; +}; + +#ifndef __SYCL_GLOBAL_VAR__ +#define __SYCL_GLOBAL_VAR__ __attribute__((sycl_global_var)) +#endif + +#define __SYCL_GLOBAL__ __attribute__((opencl_global)) + +// declaration +extern __SYCL_GLOBAL_VAR__ __SYCL_GLOBAL__ AssertHappened + __SYCL_AssertHappenedMem; + +#endif diff --git a/sycl/cmake/modules/AddSYCLUnitTest.cmake b/sycl/cmake/modules/AddSYCLUnitTest.cmake index 2d50567d2ee0b..7276927573ba9 100644 --- a/sycl/cmake/modules/AddSYCLUnitTest.cmake +++ b/sycl/cmake/modules/AddSYCLUnitTest.cmake @@ -17,10 +17,14 @@ macro(add_sycl_unittest test_dirname link_variant) if ("${link_variant}" MATCHES "SHARED") set(SYCL_LINK_LIBS ${sycl_so_target}) add_unittest(SYCLUnitTests ${test_dirname} ${ARGN}) + target_compile_definitions(${test_dirname} + PRIVATE SYCL_DISABLE_FALLBACK_ASSERT) else() add_unittest(SYCLUnitTests ${test_dirname} $ ${ARGN}) - target_compile_definitions(${test_dirname} PRIVATE __SYCL_BUILD_SYCL_DLL) + target_compile_definitions(${test_dirname} + PRIVATE __SYCL_BUILD_SYCL_DLL + SYCL_DISABLE_FALLBACK_ASSERT) get_target_property(SYCL_LINK_LIBS ${sycl_so_target} LINK_LIBRARIES) endif() @@ -66,6 +70,7 @@ macro(add_sycl_unittest_with_device test_dirname link_variant) -DGTEST_LANG_CXX11=1 -DGTEST_HAS_TR1_TUPLE=0 -D__SYCL_BUILD_SYCL_DLL + -DSYCL_DISABLE_FALLBACK_ASSERT -I${LLVM_MAIN_SRC_DIR}/utils/unittest/googletest/include -I${LLVM_MAIN_SRC_DIR}/utils/unittest/googlemock/include -I${LLVM_BINARY_DIR}/include diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index 12b074c258665..75e69a528b1c5 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -172,18 +172,25 @@ same binary image where fallback `__devicelib_assert_fail` resides. declaration: ```c++ -namespace cl { -namespace sycl { -namespace detail { -struct AssertHappened { +struct __SYCL_AssertHappened { int Flag = 0; + char Expr[256 + 1] = ""; + char File[256 + 1] = ""; + char Func[128 + 1] = ""; + + int32_t Line = 0; + + uint64_t GID0 = 0; + uint64_t GID1 = 0; + uint64_t GID2 = 0; + + uint64_t LID0 = 0; + uint64_t LID1 = 0; + uint64_t LID2 = 0; }; -} -} -} #ifdef __SYCL_DEVICE_ONLY__ -extern SYCL_GLOBAL_VAR AssertHappened AssertHappenedMem; +extern SYCL_GLOBAL_VAR __SYCL_AssertHappened __SYCL_AssertHappenedMem; #endif ``` @@ -193,6 +200,28 @@ mutable program-scope variable. The reference to extern variable is resolved within online-linking against fallback devicelib. +#### Description of fields + +The value stored here denotes if assert happened at all. There are two valid +values at host: + +| Value | Meaning | +| ----- | ------- | +| 0 | No assert failure detected | +| 2 | Assert failure detected and reported within this instance of struct | + +At device-side, there's another valid value: 1, which means that assert failure +is detected and the structure is filling up at the moment. This value is for +device-side only and should never be reported to host. Otherwise, it means, that +atomic operation malfunctioned. + +`Expr`, `File`, `Func`, `Line` are to describe the assert message itself and +contain the expression, file name, function name, line in the file where assert +failure had happened respectively. + +`GID*` and `LID*` fields describe the global and local ID respectively of a +work-item in which assert had failed. + ### Online-linking fallback `__devicelib_assert_fail` Online linking against fallback implementation of `__devicelib_assert_fail` is diff --git a/sycl/doc/PreprocessorMacros.md b/sycl/doc/PreprocessorMacros.md index ad3fd06693772..6c54547805279 100644 --- a/sycl/doc/PreprocessorMacros.md +++ b/sycl/doc/PreprocessorMacros.md @@ -33,6 +33,13 @@ SYCL 2020. Disables all deprecation warnings in SYCL runtime headers, including SYCL 1.2.1 deprecations. +### `SYCL_DISABLE_FALLBACK_ASSERT` + +Forces fallback assert feature implementation disable i.e. the *copier kernel* +and *checker host-task* are not enqueued. Also, DPCPP RT won't perform check if +user's kernel uses `__devicelib_assert_fail`. Refer to [the document](Assert.md) +for function behind *copier kernel* and *checker host-task*. + ### Version macros - `__LIBSYCL_MAJOR_VERSION` is set to SYCL runtime library major version. diff --git a/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst b/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst index b411710b773bd..d925e41b43276 100644 --- a/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst +++ b/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst @@ -14,6 +14,7 @@ cl_intel_devicelib_cassert __generic const char *func, size_t gid0, size_t gid1, size_t gid2, size_t lid0, size_t lid1, size_t lid2); + Semantic: the function is called when an assertion expression `expr` is false, and it indicates that a program does not execute as expected. @@ -21,6 +22,9 @@ The function should print a message containing the information provided in the arguments. In addition to that, the function is free to terminate the current kernel invocation. +Fallback implementation of the function raises a flag to be read later by `__devicelib_assert_read`. +The flag remains raised until the program finishes. + Arguments: - `expr` is a string representation of the assertion condition @@ -33,6 +37,16 @@ Example of a message: .. code: foo.cpp:42: void foo(int): global id: [0,0,0], local id: [0,0,0] Assertion `buf[wiID] == 0 && "Invalid value"` failed. +.. code: + int __devicelib_assert_read(); + +Semantic: +the function is called to read assert failure flag raised by +`__devicelib_assert_fail`. +The function is only used in fallback implementation. +Invoking `__devicelib_assert_fail` after a kernel doesn't imply the kernel has +assertion failed. + See also: assert_extension_. .. _assert_extension: ../Assert/SYCL_ONEAPI_ASSERT.asciidoc) diff --git a/sycl/include/CL/sycl/aspects.hpp b/sycl/include/CL/sycl/aspects.hpp index 24513ed3e515d..4ba1a4800b2dd 100644 --- a/sycl/include/CL/sycl/aspects.hpp +++ b/sycl/include/CL/sycl/aspects.hpp @@ -41,6 +41,7 @@ enum class aspect { ext_intel_mem_channel = 25, usm_atomic_host_allocations = 26, usm_atomic_shared_allocations = 27, + ext_oneapi_native_assert = 28, }; } // namespace sycl diff --git a/sycl/include/CL/sycl/detail/assert_happened.hpp b/sycl/include/CL/sycl/detail/assert_happened.hpp new file mode 100644 index 0000000000000..02ef02694cab2 --- /dev/null +++ b/sycl/include/CL/sycl/detail/assert_happened.hpp @@ -0,0 +1,40 @@ +//==------- assert_happened.hpp - Assert signalling structure --------------==// +// +// 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 + +#ifdef __SYCL_DEVICE_ONLY__ +// Reads Flag of AssertHappened on device +SYCL_EXTERNAL __attribute__((weak)) extern "C" void +__devicelib_assert_read(void *); +#endif + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace detail { +struct AssertHappened { + int Flag = 0; // set to non-zero upon assert failure + char Expr[256 + 1] = ""; + char File[256 + 1] = ""; + char Func[128 + 1] = ""; + + int32_t Line = 0; + + uint64_t GID0 = 0; + uint64_t GID1 = 0; + uint64_t GID2 = 0; + + uint64_t LID0 = 0; + uint64_t LID1 = 0; + uint64_t LID2 = 0; +}; +} // namespace detail +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/include/CL/sycl/detail/pi.def b/sycl/include/CL/sycl/detail/pi.def index 5133ed2fa9210..d5331658bec2c 100644 --- a/sycl/include/CL/sycl/detail/pi.def +++ b/sycl/include/CL/sycl/detail/pi.def @@ -133,4 +133,7 @@ _PI_API(piextPluginGetOpaqueData) _PI_API(piTearDown) +// Extension names +_PI_API(piextGetExtensionName) + #undef _PI_API diff --git a/sycl/include/CL/sycl/detail/pi.h b/sycl/include/CL/sycl/detail/pi.h index 393b5964568ff..487fa30d815d3 100644 --- a/sycl/include/CL/sycl/detail/pi.h +++ b/sycl/include/CL/sycl/detail/pi.h @@ -508,6 +508,10 @@ typedef enum { PI_PROFILING_INFO_COMMAND_END = CL_PROFILING_COMMAND_END } _pi_profiling_info; +typedef enum { + PI_INTEL_DEVICELIB_CASSERT = 0, +} _pi_extension_number; + // NOTE: this is made 64-bit to match the size of cl_mem_flags to // make the translation to OpenCL transparent. // TODO: populate @@ -572,6 +576,7 @@ using pi_program_build_info = _pi_program_build_info; using pi_program_build_status = _pi_program_build_status; using pi_kernel_info = _pi_kernel_info; using pi_profiling_info = _pi_profiling_info; +using pi_extension_number = _pi_extension_number; // For compatibility with OpenCL define this not as enum. using pi_device_partition_property = intptr_t; @@ -684,6 +689,7 @@ static const uint8_t PI_DEVICE_BINARY_OFFLOAD_KIND_SYCL = 4; #define __SYCL_PI_PROPERTY_SET_KERNEL_PARAM_OPT_INFO "SYCL/kernel param opt" /// PropertySetRegistry::SYCL_MISC_PROP defined in PropertySetIO.h #define __SYCL_PI_PROPERTY_SET_SYCL_MISC_PROP "SYCL/misc properties" +#define __SYCL_PI_PROPERTY_SET_SYCL_ASSERT_USED "SYCL/assert used" /// This struct is a record of the device binary information. If the Kind field /// denotes a portable binary type (SPIR-V or LLVM IR), the DeviceTargetSpec @@ -1649,6 +1655,19 @@ __SYCL_EXPORT pi_result piextPluginGetOpaqueData(void *opaque_data_param, /// \param PluginParameter placeholder for future use, currenly not used. __SYCL_EXPORT pi_result piTearDown(void *PluginParameter); +/// API to get extension name for those extension whose name varies with backend +/// \param ExtNumber number of extension \sa pi_extension_number +/// \param Size[out] length of returned extension name +/// \param Value[out] memory location where to put extension name +/// \return \c PI_SUCCESS if this extension name is available, +/// \c PI_INVALID_VALUE if the extension name isn't available for this +/// plugin +/// +/// Either of Size or Value isn't set if it's nil. +/// Value returned in Size can't be zero. +__SYCL_EXPORT pi_result piextGetExtensionName(pi_extension_number ExtNumber, + size_t *Size, char *Value); + struct _pi_plugin { // PI version supported by host passed to the plugin. The Plugin // checks and writes the appropriate Function Pointers in diff --git a/sycl/include/CL/sycl/detail/pi.hpp b/sycl/include/CL/sycl/detail/pi.hpp index b5e10b5a4acf3..97bb6a024e05c 100644 --- a/sycl/include/CL/sycl/detail/pi.hpp +++ b/sycl/include/CL/sycl/detail/pi.hpp @@ -338,6 +338,13 @@ class DeviceBinaryImage { const PropertyRange &getKernelParamOptInfo() const { return KernelParamOptInfo; } + const PropertyRange getAssertUsed() const { + // We can't have this variable as a class member, since it would break + // the ABI backwards compatibility. + PropertyRange AssertUsed; + AssertUsed.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_ASSERT_USED); + return AssertUsed; + } virtual ~DeviceBinaryImage() {} protected: diff --git a/sycl/include/CL/sycl/event.hpp b/sycl/include/CL/sycl/event.hpp index 41dc3bb557597..0a9ca3aba820e 100644 --- a/sycl/include/CL/sycl/event.hpp +++ b/sycl/include/CL/sycl/event.hpp @@ -133,7 +133,6 @@ class __SYCL_EXPORT event { return reinterpret_cast::type>( getNative()); } - private: event(shared_ptr_class EventImpl); diff --git a/sycl/include/CL/sycl/handler.hpp b/sycl/include/CL/sycl/handler.hpp index ad9b003c51a3f..f26b043aa8e6f 100644 --- a/sycl/include/CL/sycl/handler.hpp +++ b/sycl/include/CL/sycl/handler.hpp @@ -2322,6 +2322,7 @@ class __SYCL_EXPORT handler { access::target); friend class ::MockHandler; + friend class detail::queue_impl; template auto getRangeRoundedKernelLambda(KernelType KernelFunc, diff --git a/sycl/include/CL/sycl/queue.hpp b/sycl/include/CL/sycl/queue.hpp index 0dfcb7ecbb98f..1de59de3a9262 100644 --- a/sycl/include/CL/sycl/queue.hpp +++ b/sycl/include/CL/sycl/queue.hpp @@ -9,6 +9,7 @@ #pragma once #include +#include #include #include #include @@ -20,6 +21,7 @@ #include #include +#include #include // having _TWO_ mid-param #ifdefs makes the functions very difficult to read. @@ -58,14 +60,29 @@ #define _KERNELFUNCPARAM(a) const KernelType &a #endif +// Helper macro to identify if fallback assert is needed +// FIXME remove __NVPTX__ condition once devicelib supports CUDA +#if !defined(SYCL_DISABLE_FALLBACK_ASSERT) && !defined(__NVPTX__) +#define __SYCL_USE_FALLBACK_ASSERT 1 +#else +#define __SYCL_USE_FALLBACK_ASSERT 0 +#endif + __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { // Forward declaration class context; class device; +class queue; + namespace detail { class queue_impl; +#if __SYCL_USE_FALLBACK_ASSERT +class AssertInfoCopier; +static event submitAssertCapture(queue &, event &, queue *, + const detail::code_location &); +#endif } /// Encapsulates a single SYCL queue which schedules kernels on a SYCL device. @@ -214,6 +231,7 @@ class __SYCL_EXPORT queue { template typename info::param_traits::return_type get_info() const; +public: /// Submits a command group function object to the queue, in order to be /// scheduled for execution on the device. /// @@ -223,7 +241,30 @@ class __SYCL_EXPORT queue { template event submit(T CGF _CODELOCPARAM(&CodeLoc)) { _CODELOCARG(&CodeLoc); - return submit_impl(CGF, CodeLoc); + event Event; + +#if __SYCL_USE_FALLBACK_ASSERT + if (!is_host()) { + auto PostProcess = [this, &CodeLoc](bool IsKernel, bool KernelUsesAssert, + event &E) { + if (IsKernel && !get_device().has(aspect::ext_oneapi_native_assert) && + KernelUsesAssert) { + // __devicelib_assert_fail isn't supported by Device-side Runtime + // Linking against fallback impl of __devicelib_assert_fail is + // performed by program manager class + submitAssertCapture(*this, E, /* SecondaryQueue = */ nullptr, + CodeLoc); + } + }; + + Event = submit_impl_and_postprocess(CGF, CodeLoc, PostProcess); + } else +#endif // !defined(SYCL_DISABLE_FALLBACK_ASSERT) && !defined(__NVPTX__) + { + Event = submit_impl(CGF, CodeLoc); + } + + return Event; } /// Submits a command group function object to the queue, in order to be @@ -241,7 +282,27 @@ class __SYCL_EXPORT queue { event submit(T CGF, queue &SecondaryQueue _CODELOCPARAM(&CodeLoc)) { _CODELOCARG(&CodeLoc); - return submit_impl(CGF, SecondaryQueue, CodeLoc); + event Event; + +#if __SYCL_USE_FALLBACK_ASSERT + auto PostProcess = [this, &SecondaryQueue, &CodeLoc]( + bool IsKernel, bool KernelUsesAssert, event &E) { + if (IsKernel && !get_device().has(aspect::ext_oneapi_native_assert) && + KernelUsesAssert) { + // __devicelib_assert_fail isn't supported by Device-side Runtime + // Linking against fallback impl of __devicelib_assert_fail is performed + // by program manager class + submitAssertCapture(*this, E, /* SecondaryQueue = */ nullptr, CodeLoc); + } + }; + + Event = + submit_impl_and_postprocess(CGF, SecondaryQueue, CodeLoc, PostProcess); +#else + Event = submit_impl(CGF, SecondaryQueue, CodeLoc); +#endif // !defined(SYCL_DISABLE_FALLBACK_ASSERT) && !defined(__NVPTX__) + + return Event; } /// Prevents any commands submitted afterward to this queue from executing @@ -938,6 +999,11 @@ class __SYCL_EXPORT queue { template friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj); +#if __SYCL_USE_FALLBACK_ASSERT + friend event detail::submitAssertCapture(queue &, event &, queue *, + const detail::code_location &); +#endif + /// A template-free version of submit. event submit_impl(function_class CGH, const detail::code_location &CodeLoc); @@ -945,6 +1011,33 @@ class __SYCL_EXPORT queue { event submit_impl(function_class CGH, queue secondQueue, const detail::code_location &CodeLoc); + // Function to postprocess submitted command + // Arguments: + // bool IsKernel - true if the submitted command was kernel, false otherwise + // bool KernelUsesAssert - true if submitted kernel uses assert, only + // meaningful when IsKernel is true + // event &Event - event after which post processing should be executed + using SubmitPostProcessF = std::function; + + /// A template-free version of submit. + /// \param CGH command group function/handler + /// \param CodeLoc code location + /// + /// This method stores additional information within event_impl class instance + event submit_impl_and_postprocess(function_class CGH, + const detail::code_location &CodeLoc, + const SubmitPostProcessF &PostProcess); + /// A template-free version of submit. + /// \param CGH command group function/handler + /// \param secondQueue fallback queue + /// \param CodeLoc code location + /// + /// This method stores additional information within event_impl class instance + event submit_impl_and_postprocess(function_class CGH, + queue secondQueue, + const detail::code_location &CodeLoc, + const SubmitPostProcessF &PostProcess); + /// parallel_for_impl with a kernel represented as a lambda + range that /// specifies global size only. /// @@ -1007,8 +1100,93 @@ class __SYCL_EXPORT queue { }, CodeLoc); } + + buffer &getAssertHappenedBuffer(); }; +namespace detail { +#if __SYCL_USE_FALLBACK_ASSERT +#define __SYCL_ASSERT_START 1 +/** + * Submit copy task for assert failure flag and host-task to check the flag + * \param Event kernel's event to depend on i.e. the event represents the + * kernel to check for assertion failure + * \param SecondaryQueue secondary queue for submit process, null if not used + * \returns host tasks event + * + * This method doesn't belong to queue class to overcome msvc behaviour due to + * which it gets compiled and exported without any integration header and, thus, + * with no proper KernelInfo instance. + */ +event submitAssertCapture(queue &Self, event &Event, queue *SecondaryQueue, + const detail::code_location &CodeLoc) { + using AHBufT = buffer; + + AHBufT &Buffer = Self.getAssertHappenedBuffer(); + + event CopierEv, CheckerEv, PostCheckerEv; + auto CopierCGF = [&](handler &CGH) { + CGH.depends_on(Event); + + auto Acc = Buffer.get_access(CGH); + + CGH.single_task([Acc] { +#ifdef __SYCL_DEVICE_ONLY__ + __devicelib_assert_read(&Acc[0]); +#else + (void)Acc; +#endif // __SYCL_DEVICE_ONLY__ + }); + }; + auto CheckerCGF = [&CopierEv, &Buffer](handler &CGH) { + CGH.depends_on(CopierEv); + using mode = access::mode; + using target = access::target; + + auto Acc = Buffer.get_access(CGH); + + CGH.codeplay_host_task([=] { + const detail::AssertHappened *AH = &Acc[0]; + + // Don't use assert here as msvc will insert reference to __imp__wassert + // which won't be properly resolved in separate compile use-case +#ifndef NDEBUG + if (AH->Flag == __SYCL_ASSERT_START) + throw sycl::runtime_error( + "Internal Error. Invalid value in assert description.", + PI_INVALID_VALUE); +#endif + + if (AH->Flag) { + const char *Expr = AH->Expr[0] ? AH->Expr : ""; + const char *File = AH->File[0] ? AH->File : ""; + const char *Func = AH->Func[0] ? AH->Func : ""; + + fprintf(stderr, + "%s:%d: %s: global id: [%" PRIu64 ",%" PRIu64 ",%" PRIu64 + "], local id: [%" PRIu64 ",%" PRIu64 ",%" PRIu64 "] " + "Assertion `%s` failed.\n", + File, AH->Line, Func, AH->GID0, AH->GID1, AH->GID2, AH->LID0, + AH->LID1, AH->LID2, Expr); + abort(); // no need to release memory as it's abort anyway + } + }); + }; + + if (SecondaryQueue) { + CopierEv = Self.submit_impl(CopierCGF, *SecondaryQueue, CodeLoc); + CheckerEv = Self.submit_impl(CheckerCGF, *SecondaryQueue, CodeLoc); + } else { + CopierEv = Self.submit_impl(CopierCGF, CodeLoc); + CheckerEv = Self.submit_impl(CheckerCGF, CodeLoc); + } + + return CheckerEv; +} +#undef __SYCL_ASSERT_START +#endif // !defined(SYCL_DISABLE_FALLBACK_ASSERT) && !defined(__NVPTX__) +} // namespace detail + } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) @@ -1021,3 +1199,5 @@ template <> struct hash { } }; } // namespace std + +#undef __SYCL_USE_FALLBACK_ASSERT diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index 1c8daf6e4b835..dc37fce914d29 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -4629,6 +4629,27 @@ pi_result cuda_piextUSMGetMemAllocInfo(pi_context context, const void *ptr, // pi_level_zero.cpp for reference) Currently this is just a NOOP. pi_result cuda_piTearDown(void *) { return PI_SUCCESS; } +pi_result cuda_piextGetExtensionName(pi_extension_number ExtNumber, + size_t *Size, char *Value) { + pi_result Result = PI_SUCCESS; + // TODO switch to map/unordered_map when have enough number of extensions + switch (ExtNumber) { + case PI_INTEL_DEVICELIB_CASSERT: { + // FIXME set name after backend support + static const std::string Name = "N/A"; + if (Size) + *Size = Name.length(); + if (Value) + std::memcpy(Value, Name.data(), Name.length()); + break; + } + default: + Result = PI_INVALID_VALUE; + } + + return Result; +} + const char SupportedVersion[] = _PI_H_VERSION_STRING; pi_result piPluginInit(pi_plugin *PluginInit) { @@ -4769,6 +4790,8 @@ pi_result piPluginInit(pi_plugin *PluginInit) { _PI_CL(piextKernelSetArgSampler, cuda_piextKernelSetArgSampler) _PI_CL(piTearDown, cuda_piTearDown) + _PI_CL(piextGetExtensionName, cuda_piextGetExtensionName); + #undef _PI_CL return PI_SUCCESS; diff --git a/sycl/plugins/esimd_cpu/pi_esimd_cpu.cpp b/sycl/plugins/esimd_cpu/pi_esimd_cpu.cpp index 63fc720f49eee..9a9fe2e8f8f57 100644 --- a/sycl/plugins/esimd_cpu/pi_esimd_cpu.cpp +++ b/sycl/plugins/esimd_cpu/pi_esimd_cpu.cpp @@ -1232,6 +1232,27 @@ pi_result piTearDown(void *) { return PI_SUCCESS; } +pi_result piextGetExtensionName(pi_extension_number ExtNumber, size_t *Size, + char *Value) { + pi_result Result = PI_SUCCESS; + // TODO switch to map/unordered_map when have enough number of extensions + switch (ExtNumber) { + case PI_INTEL_DEVICELIB_CASSERT: { + // FIXME set name after backend support + static const std::string Name = "N/A"; + if (Size) + *Size = Name.length(); + if (Value) + std::memcpy(Value, Name.data(), Name.length()); + break; + } + default: + Result = PI_INVALID_VALUE; + } + + return Result; +} + pi_result piPluginInit(pi_plugin *PluginInit) { assert(PluginInit); size_t PluginVersionSize = sizeof(PluginInit->PluginVersion); diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index a22b1a729dfc8..d9782121b4d2c 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -6545,4 +6545,25 @@ pi_result piTearDown(void *PluginParameter) { return PI_SUCCESS; } +pi_result piextGetExtensionName(pi_extension_number ExtNumber, size_t *Size, + char *Value) { + pi_result Result = PI_SUCCESS; + // TODO switch to map/unordered_map when have enough number of extensions + switch (ExtNumber) { + case PI_INTEL_DEVICELIB_CASSERT: { + // FIXME set name after backend support + static const std::string Name = "N/A"; + if (Size) + *Size = Name.length(); + if (Value) + std::memcpy(Value, Name.data(), Name.length()); + break; + } + default: + Result = PI_INVALID_VALUE; + } + + return Result; +} + } // extern "C" diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index f1a074b3fc65d..b9f5645bb2c79 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -1208,6 +1208,26 @@ pi_result piTearDown(void *PluginParameter) { return PI_SUCCESS; } +pi_result piextGetExtensionName(pi_extension_number ExtNumber, size_t *Size, + char *Value) { + pi_result Result = PI_SUCCESS; + // TODO switch to map/unordered_map when have enough number of extensions + switch (ExtNumber) { + case PI_INTEL_DEVICELIB_CASSERT: { + static const std::string Name = "cl_intel_devicelib_cassert"; + if (Size) + *Size = Name.length(); + if (Value) + std::memcpy(Value, Name.data(), Name.length()); + break; + } + default: + Result = PI_INVALID_VALUE; + } + + return Result; +} + pi_result piPluginInit(pi_plugin *PluginInit) { int CompareVersions = strcmp(PluginInit->PiVersion, SupportedVersion); if (CompareVersions < 0) { @@ -1339,6 +1359,8 @@ pi_result piPluginInit(pi_plugin *PluginInit) { _PI_CL(piextKernelSetArgSampler, piextKernelSetArgSampler) _PI_CL(piTearDown, piTearDown) + _PI_CL(piextGetExtensionName, piextGetExtensionName); + #undef _PI_CL return PI_SUCCESS; diff --git a/sycl/source/detail/device_impl.cpp b/sycl/source/detail/device_impl.cpp index 172c19be83344..d83d87378ce0b 100644 --- a/sycl/source/detail/device_impl.cpp +++ b/sycl/source/detail/device_impl.cpp @@ -300,6 +300,8 @@ bool device_impl::has(aspect Aspect) const { case aspect::ext_intel_max_mem_bandwidth: // currently not supported return false; + case aspect::ext_oneapi_native_assert: + return isAssertFailSupported(); default: throw runtime_error("This device aspect has not been implemented yet.", @@ -314,6 +316,30 @@ std::shared_ptr device_impl::getHostDeviceImpl() { return HostImpl; } +bool device_impl::isAssertFailSupported() const { + // assert is sort of natively supported by host + if (MIsHostDevice) + return true; + + const plugin &Plugin = getPlugin(); + + if (Plugin.getBackend() == backend::cuda) + return true; + + size_t ExtNameSize = 0; + Plugin.call(PI_INTEL_DEVICELIB_CASSERT, + &ExtNameSize, nullptr); + + assert(ExtNameSize && "Size can't be zero"); + + std::unique_ptr ExtName{new char[ExtNameSize + 1]}; + ExtName[ExtNameSize] = '\0'; + Plugin.call(PI_INTEL_DEVICELIB_CASSERT, + nullptr, ExtName.get()); + + return has_extension(ExtName.get()); +} + } // namespace detail } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/source/detail/device_impl.hpp b/sycl/source/detail/device_impl.hpp index 0e1381f933964..93e1d963faf40 100644 --- a/sycl/source/detail/device_impl.hpp +++ b/sycl/source/detail/device_impl.hpp @@ -222,6 +222,8 @@ class device_impl { /// \return the host device_impl singleton static std::shared_ptr getHostDeviceImpl(); + bool isAssertFailSupported() const; + private: explicit device_impl(pi_native_handle InteropDevice, RT::PiDevice Device, PlatformImplPtr Platform, const plugin &Plugin); diff --git a/sycl/source/detail/queue_impl.cpp b/sycl/source/detail/queue_impl.cpp index 73f1943d638e3..cf59a6f52b6e3 100644 --- a/sycl/source/detail/queue_impl.cpp +++ b/sycl/source/detail/queue_impl.cpp @@ -287,6 +287,20 @@ pi_native_handle queue_impl::getNative() const { return Handle; } +bool queue_impl::kernelUsesAssert(const std::string &KernelName, + OSModuleHandle Handle) const { + RTDeviceBinaryImage &BinImg = ProgramManager::getInstance().getDeviceImage( + Handle, KernelName, get_context(), get_device()); + + const pi::DeviceBinaryImage::PropertyRange &AssertUsedRange = + BinImg.getAssertUsed(); + if (AssertUsedRange.isAvailable()) + for (const auto &Prop : AssertUsedRange) + if (Prop->Name == KernelName) + return true; + + return false; +} } // namespace detail } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/source/detail/queue_impl.hpp b/sycl/source/detail/queue_impl.hpp index 1690786eebfec..e5f69b0164406 100644 --- a/sycl/source/detail/queue_impl.hpp +++ b/sycl/source/detail/queue_impl.hpp @@ -9,6 +9,7 @@ #pragma once #include +#include #include #include #include @@ -78,7 +79,8 @@ class queue_impl { queue_impl(const DeviceImplPtr &Device, const ContextImplPtr &Context, const async_handler &AsyncHandler, const property_list &PropList) : MDevice(Device), MContext(Context), MAsyncHandler(AsyncHandler), - MPropList(PropList), MHostQueue(MDevice->is_host()) { + MPropList(PropList), MHostQueue(MDevice->is_host()), + MAssertHappenedBuffer(range<1>{1}) { if (!Context->hasDevice(Device)) throw cl::sycl::invalid_parameter_error( "Queue cannot be constructed with the given context and device " @@ -101,7 +103,8 @@ class queue_impl { /// \param AsyncHandler is a SYCL asynchronous exception handler. queue_impl(RT::PiQueue PiQueue, const ContextImplPtr &Context, const async_handler &AsyncHandler) - : MContext(Context), MAsyncHandler(AsyncHandler), MHostQueue(false) { + : MContext(Context), MAsyncHandler(AsyncHandler), MHostQueue(false), + MAssertHappenedBuffer(range<1>{1}) { MQueues.push_back(pi::cast(PiQueue)); @@ -156,6 +159,8 @@ class queue_impl { template typename info::param_traits::return_type get_info() const; + using SubmitPostProcessF = std::function; + /// Submits a command group function object to the queue, in order to be /// scheduled for execution on the device. /// @@ -166,20 +171,22 @@ class queue_impl { /// \param Self is a shared_ptr to this queue. /// \param SecondQueue is a shared_ptr to the secondary queue. /// \param Loc is the code location of the submit call (default argument) + /// \param StoreAdditionalInfo makes additional info be stored in event_impl /// \return a SYCL event object, which corresponds to the queue the command /// group is being enqueued on. event submit(const function_class &CGF, const shared_ptr_class &Self, const shared_ptr_class &SecondQueue, - const detail::code_location &Loc) { + const detail::code_location &Loc, + const SubmitPostProcessF *PostProcess = nullptr) { try { - return submit_impl(CGF, Self, Loc); + return submit_impl(CGF, Self, Loc, PostProcess); } catch (...) { { std::lock_guard Lock(MMutex); MExceptions.PushBack(std::current_exception()); } - return SecondQueue->submit(CGF, SecondQueue, Loc); + return SecondQueue->submit(CGF, SecondQueue, Loc, PostProcess); } } @@ -189,11 +196,13 @@ class queue_impl { /// \param CGF is a function object containing command group. /// \param Self is a shared_ptr to this queue. /// \param Loc is the code location of the submit call (default argument) + /// \param StoreAdditionalInfo makes additional info be stored in event_impl /// \return a SYCL event object for the submitted command group. event submit(const function_class &CGF, const shared_ptr_class &Self, - const detail::code_location &Loc) { - return submit_impl(CGF, Self, Loc); + const detail::code_location &Loc, + const SubmitPostProcessF *PostProcess = nullptr) { + return submit_impl(CGF, Self, Loc, PostProcess); } /// Performs a blocking wait for the completion of all enqueued tasks in the @@ -387,20 +396,46 @@ class queue_impl { /// \return a native handle. pi_native_handle getNative() const; + bool kernelUsesAssert(const std::string &KernelName, + OSModuleHandle Handle) const; + + void asynchronouslyDeleteBuffer(buffer *B); + + buffer &getAssertHappenedBuffer() { + return MAssertHappenedBuffer; + } + private: /// Performs command group submission to the queue. /// /// \param CGF is a function object containing command group. + /// \param[out] IsKernel set to true if kernel was submit /// \param Self is a pointer to this queue. /// \param Loc is the code location of the submit call (default argument) /// \return a SYCL event representing submitted command group. event submit_impl(const function_class &CGF, const shared_ptr_class &Self, - const detail::code_location &Loc) { + const detail::code_location &Loc, + const SubmitPostProcessF *PostProcess) { handler Handler(Self, MHostQueue); Handler.saveCodeLoc(Loc); CGF(Handler); - event Event = Handler.finalize(); + + event Event; + + if (PostProcess) { + bool IsKernel = Handler.getType() == CG::KERNEL; + bool KernelUsesAssert = false; + if (IsKernel) + KernelUsesAssert = + kernelUsesAssert(Handler.MKernelName, Handler.MOSModuleHandle); + + Event = Handler.finalize(); + + (*PostProcess)(IsKernel, KernelUsesAssert, Event); + } else + Event = Handler.finalize(); + addEvent(Event); return Event; } @@ -457,6 +492,9 @@ class queue_impl { // Thread pool for host task and event callbacks execution. // The thread pool is instantiated upon the very first call to getThreadPool() std::unique_ptr MHostTaskThreadPool; + + // Buffer to store assert failure descriptor + buffer MAssertHappenedBuffer; }; } // namespace detail diff --git a/sycl/source/detail/scheduler/scheduler.hpp b/sycl/source/detail/scheduler/scheduler.hpp index 5193d61858849..bae37a5d5e88f 100644 --- a/sycl/source/detail/scheduler/scheduler.hpp +++ b/sycl/source/detail/scheduler/scheduler.hpp @@ -763,6 +763,7 @@ class Scheduler { friend class Command; friend class DispatchHostTask; + friend class queue_impl; /// Stream buffers structure. /// diff --git a/sycl/source/event.cpp b/sycl/source/event.cpp index c8b94a5644b13..2c73af98b2c38 100644 --- a/sycl/source/event.cpp +++ b/sycl/source/event.cpp @@ -89,6 +89,5 @@ event::event(shared_ptr_class event_impl) backend event::get_backend() const noexcept { return getImplBackend(impl); } pi_native_handle event::getNative() const { return impl->getNative(); } - } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/source/queue.cpp b/sycl/source/queue.cpp index 5a0209d7f6fc0..9e8fb721812f5 100644 --- a/sycl/source/queue.cpp +++ b/sycl/source/queue.cpp @@ -128,6 +128,19 @@ event queue::submit_impl(function_class CGH, queue SecondQueue, return impl->submit(CGH, impl, SecondQueue.impl, CodeLoc); } +event queue::submit_impl_and_postprocess( + function_class CGH, const detail::code_location &CodeLoc, + const SubmitPostProcessF &PostProcess) { + return impl->submit(CGH, impl, CodeLoc, &PostProcess); +} + +event queue::submit_impl_and_postprocess( + function_class CGH, queue SecondQueue, + const detail::code_location &CodeLoc, + const SubmitPostProcessF &PostProcess) { + return impl->submit(CGH, impl, SecondQueue.impl, CodeLoc, &PostProcess); +} + void queue::wait_proxy(const detail::code_location &CodeLoc) { impl->wait(CodeLoc); } @@ -171,5 +184,8 @@ backend queue::get_backend() const noexcept { return getImplBackend(impl); } pi_native_handle queue::getNative() const { return impl->getNative(); } +buffer &queue::getAssertHappenedBuffer() { + return impl->getAssertHappenedBuffer(); +} } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/test/abi/pi_level_zero_symbol_check.dump b/sycl/test/abi/pi_level_zero_symbol_check.dump index f080521418500..2a0b7f090f95a 100644 --- a/sycl/test/abi/pi_level_zero_symbol_check.dump +++ b/sycl/test/abi/pi_level_zero_symbol_check.dump @@ -88,6 +88,7 @@ piextDeviceSelectBinary piextEventCreateWithNativeHandle piextEventGetNativeHandle piextGetDeviceFunctionPointer +piextGetExtensionName piextKernelCreateWithNativeHandle piextKernelGetNativeHandle piextKernelSetArgMemObj diff --git a/sycl/test/abi/pi_opencl_symbol_check.dump b/sycl/test/abi/pi_opencl_symbol_check.dump index 23f7a3992bf60..d3d2dabfb2aa0 100644 --- a/sycl/test/abi/pi_opencl_symbol_check.dump +++ b/sycl/test/abi/pi_opencl_symbol_check.dump @@ -34,6 +34,7 @@ piextDeviceGetNativeHandle piextDeviceSelectBinary piextEventCreateWithNativeHandle piextGetDeviceFunctionPointer +piextGetExtensionName piextKernelCreateWithNativeHandle piextKernelGetNativeHandle piextKernelSetArgMemObj diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 2016f1ae315c0..98f5c66f8d282 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3654,6 +3654,9 @@ _ZN2cl4sycl5queue11submit_implESt8functionIFvRNS0_7handlerEEERKNS0_6detail13code _ZN2cl4sycl5queue11submit_implESt8functionIFvRNS0_7handlerEEES1_RKNS0_6detail13code_locationE _ZN2cl4sycl5queue18throw_asynchronousEv _ZN2cl4sycl5queue20wait_and_throw_proxyERKNS0_6detail13code_locationE +_ZN2cl4sycl5queue23getAssertHappenedBufferEv +_ZN2cl4sycl5queue27submit_impl_and_postprocessESt8functionIFvRNS0_7handlerEEERKNS0_6detail13code_locationERKS2_IFvbbRNS0_5eventEEE +_ZN2cl4sycl5queue27submit_impl_and_postprocessESt8functionIFvRNS0_7handlerEEES1_RKNS0_6detail13code_locationERKS2_IFvbbRNS0_5eventEEE _ZN2cl4sycl5queue6memcpyEPvPKvm _ZN2cl4sycl5queue6memcpyEPvPKvmNS0_5eventE _ZN2cl4sycl5queue6memcpyEPvPKvmRKSt6vectorINS0_5eventESaIS6_EE diff --git a/sycl/test/on-device/plugins/level_zero_dynamic_batch_test.cpp b/sycl/test/on-device/plugins/level_zero_dynamic_batch_test.cpp index bc07034e05110..3dc8c3c7da102 100644 --- a/sycl/test/on-device/plugins/level_zero_dynamic_batch_test.cpp +++ b/sycl/test/on-device/plugins/level_zero_dynamic_batch_test.cpp @@ -1,6 +1,7 @@ // REQUIRES: gpu, level_zero - -// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// Disable fallback assert as it employs host-task which calls event::wait and +// causes to execute open command lists +// RUN: %clangxx -DSYCL_DISABLE_FALLBACK_ASSERT -fsycl -fsycl-targets=%sycl_triple %s -o %t.out // Check that dynamic batching raises/lowers batch size // RUN: env SYCL_PI_TRACE=2 ZE_DEBUG=1 %GPU_RUN_PLACEHOLDER %t.out 2>&1 | FileCheck --check-prefixes=CKALL,CKDYN %s