From 2911ea78aa52699334d98c1066e20200a6e5000b Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 31 Mar 2021 17:07:13 +0300 Subject: [PATCH 01/45] [SYCL] [DOC] Prepare design-document for assert feature Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 144 +++++++++++++++++++++++ sycl/doc/extensions/Assert/abort.md | 1 + sycl/doc/extensions/Assert/level-zero.md | 19 +++ sycl/doc/extensions/Assert/opencl.md | 22 ++++ 4 files changed, 186 insertions(+) create mode 100644 sycl/doc/Assert.md create mode 100644 sycl/doc/extensions/Assert/abort.md create mode 100644 sycl/doc/extensions/Assert/level-zero.md create mode 100644 sycl/doc/extensions/Assert/opencl.md diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md new file mode 100644 index 0000000000000..d4582d861c5c3 --- /dev/null +++ b/sycl/doc/Assert.md @@ -0,0 +1,144 @@ +# Assert feature + +**IMPORTANT**: This document is a draft. + +During debugging of kernel code user may put assertions here and there. +The expected behaviour of assertion failure at host is application abort. +Our choice for device-side assertions is asynchronous exception in order to +allow for extensibility. + +The user is free to disable assertions by defining `NDEBUG` macro at +compile-time. + + +## Use-case example + +``` +using namespace cl::sycl; +auto ErrorHandler = [] (exception_list Exs) { + for (exception_ptr const& E : Exs) { + try { + std::rethrow_exception(E); + } + catch (event_error const& Ex) { + std::cout << “Exception - ” << Ex.what(); // assertion failed + std::abort(); + } + } +}; + +void user_func(item<2> Item) { + assert((Item[0] % 2) && “Nil”); +} + +int main() { + queue Q(ErrorHandler); + q.submit([&] (handler& CGH) { + CGH.parallel_for(range<2>{N, M}, [=](item<2> It) { + do_smth(); + user_func(It); + do_smth_else(); + }); + }); + Q.wait_and_throw(); + std::cout << “One shouldn’t see this message.“; + return 0; +} +``` + +In this use-case every work-item with even X dimension will trigger assertion +failure. Assertion failure should be reported via asynchronous exceptions. If +asynchronous exception handler is set the failure is reported with +`cl::sycl::event_error` exception. Otherwise, SYCL Runtime should trigger abort. +At least one failed assertion should be reported. + +When multiple kernels are enqueued and both fail at assertion at least single +assertion should be reported. + +## User requirements + +From user's point of view there are the following requirements: + +| # | Title | Description | Importance | +| - | ----- | ----------- | ---------- | +| 1 | Handle assertion failure | Signal about assertion failure via SYCL asynchronous exception | Must have | +| 2 | Print assert message | Assert function should print message to stderr at host | Must have | +| 3 | Stop under debugger | When debugger is attached, break at assertion point | Highly desired | +| 4 | Reliability | Assert failure should be reported regardless of kernel deadlock | Highly desired | + +## Contents of `cl::sycl::event_error` + +`cl::sycl::event_error::what()` should return the same assertion failure message +as is printed at the time being. + +Other than that, interface of `cl::sycl::event_error` should look like: +``` +class event_error : public runtime_error { +public: + event_error() = default; + + event_error(const char *Msg, cl_int Err) + : event_error(string_class(Msg), Err) {} + + event_error(const string_class &Msg, cl_int Err) : runtime_error(Msg, Err) {} + + /// Returns global ID with the dimension provided + int globalId(int Dim) const; + + /// Returns local ID with the dimension provided + int localId(int Dim) const; +}; +``` + +Regardless of whether asynchronous exception handler is set or not, there's an +action to be performed by SYCL Runtime. To achieve this, information about +assert failure should be propagated from device-side to SYCL Runtime. This +should be performed via calls to `clGetEventInfo` for OpenCL backend and +`zeEventQueryStatus` for Level-Zero backend. + +## Terms + + - Device-side Runtime - part of device-code, which is supplied by Device-side + Compiler. + - Low-level Runtime - the backend/runtime, behind DPCPP Runtime. + - Device-side Compiler - compiler which generates device-native bitcode based + on input SPIR-V image. + - Accessor metadata - parts of accessor representation at device-side: pointer, + ranges, offset. + +## How it works? + +For the time being, `assert(expr)` macro ends up in call to +`__devicelib_assert_fail` function. This function is part of [Device library extension](doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst#cl_intel_devicelib_cassert). +Device code already contains call to the function. Currently, a device-binary +is always linked against fallback implementation. +Device-side compiler/linker provides their implementation of `__devicelib_assert_fail` +and prefer this implementation over fallback one. + +If Device-side Runtime supports `__devicelib_assert_fail` then Low-Level Runtime +is responsible for: + - detecting if assert failure took place; + - flushing assert message to `stderr` on host. +When detected, Low-level Runtime reports assert failure to DPCPP Runtime +at synchronization points. + +Refer to [OpenCL](doc/extensions/Assert/opencl.md) and [Level-Zero](doc/extensions/Assert/level-zero.md) +extensions. + +If Device-side Runtime doesn't support `__devicelib_assert_fail` then a buffer +based approach comes in place. The approach doesn't require any support from +Device-side Runtime. Neither it does from Low-level Runtime. + +Within this approach, a dedicated assert buffer is allocated and implicit kernel +argument is introduced. The argument is an accessor with `discard_read_write` +or `discard_write` access mode. Accessor metadata is stored to program scope +variable. This allows to refer to the accessor without modifying each and every +user's function. Fallback implementation of `__devicelib_assert_fail` restores +accessor metadata from program scope variable and writes assert information to +the assert buffer. Atomic operations are used in order to not overwrite existing +information. + +Storing and restoring of accessor metadata to/from program scope variable is +performed with help of builtins. Implementations of these builtins are +substituted by frontend. + diff --git a/sycl/doc/extensions/Assert/abort.md b/sycl/doc/extensions/Assert/abort.md new file mode 100644 index 0000000000000..8b137891791fe --- /dev/null +++ b/sycl/doc/extensions/Assert/abort.md @@ -0,0 +1 @@ + diff --git a/sycl/doc/extensions/Assert/level-zero.md b/sycl/doc/extensions/Assert/level-zero.md new file mode 100644 index 0000000000000..265ff5858d83f --- /dev/null +++ b/sycl/doc/extensions/Assert/level-zero.md @@ -0,0 +1,19 @@ +# Overview + +This extension enables detection of assert failure of kernel. + +# New enum value + +`ze_result_t` enumeration should be augmented with `ZE_RESULT_ABORTED` enum +element. This enum value indicated a detected assert failure at device-side. + +# Changed API + +``` +ze_event_handle_t Event; // describes an event of kernel been submitted previously +ze_result Result = zeEventQueryStatus(Event); +``` + +If kernel failed an assertion `zeEventQueryStatus` should return +`ZE_RESULT_ABORTED`. + diff --git a/sycl/doc/extensions/Assert/opencl.md b/sycl/doc/extensions/Assert/opencl.md new file mode 100644 index 0000000000000..50ad0b7db0897 --- /dev/null +++ b/sycl/doc/extensions/Assert/opencl.md @@ -0,0 +1,22 @@ +# Overview + +This extension enables detection of assert failure of kernel. + +# New error code + +`CL_ASSERT_FAILURE` is added to indicate a detected assert failure at +device-side. + +# Changed API + +``` +cl_event Event; // describes an event of kernel been submitted previously +cl_int Result; +size_t ResultSize; + +clGetEventInfo(Event, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(Result), &Result, &ResultSize); +``` + +If kernel failed an assertion `clGetEventInfo` should put `CL_ASSERT_FAILURE` +in `Result`. + From b69a1cdf1f99b5ef5c78a9d0db0fc92ddccfd64a Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 31 Mar 2021 17:43:38 +0300 Subject: [PATCH 02/45] Remove redundant file Signed-off-by: Sergey Kanaev --- sycl/doc/extensions/Assert/abort.md | 1 - 1 file changed, 1 deletion(-) delete mode 100644 sycl/doc/extensions/Assert/abort.md diff --git a/sycl/doc/extensions/Assert/abort.md b/sycl/doc/extensions/Assert/abort.md deleted file mode 100644 index 8b137891791fe..0000000000000 --- a/sycl/doc/extensions/Assert/abort.md +++ /dev/null @@ -1 +0,0 @@ - From 15ea88ea0a43595e33245a5852559794a3b52eee Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Thu, 1 Apr 2021 13:57:14 +0300 Subject: [PATCH 03/45] Fix typo Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 6 ------ 1 file changed, 6 deletions(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index d4582d861c5c3..3b5cbfbfafe47 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -81,12 +81,6 @@ public: : event_error(string_class(Msg), Err) {} event_error(const string_class &Msg, cl_int Err) : runtime_error(Msg, Err) {} - - /// Returns global ID with the dimension provided - int globalId(int Dim) const; - - /// Returns local ID with the dimension provided - int localId(int Dim) const; }; ``` From ca08fecc0945c673579505bf34eee6abba74b902 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Mon, 5 Apr 2021 16:16:43 +0300 Subject: [PATCH 04/45] Address some review comments. Add description of built-ins. Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 130 +++++++++++++++++++++++++++++++++++++++++---- 1 file changed, 119 insertions(+), 11 deletions(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index 3b5cbfbfafe47..956fd64732cf5 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -55,6 +55,7 @@ At least one failed assertion should be reported. When multiple kernels are enqueued and both fail at assertion at least single assertion should be reported. + ## User requirements From user's point of view there are the following requirements: @@ -66,12 +67,10 @@ From user's point of view there are the following requirements: | 3 | Stop under debugger | When debugger is attached, break at assertion point | Highly desired | | 4 | Reliability | Assert failure should be reported regardless of kernel deadlock | Highly desired | -## Contents of `cl::sycl::event_error` -`cl::sycl::event_error::what()` should return the same assertion failure message -as is printed at the time being. +## Contents of `cl::sycl::event_error` -Other than that, interface of `cl::sycl::event_error` should look like: +Interface of `cl::sycl::event_error` should look like: ``` class event_error : public runtime_error { public: @@ -90,22 +89,28 @@ assert failure should be propagated from device-side to SYCL Runtime. This should be performed via calls to `clGetEventInfo` for OpenCL backend and `zeEventQueryStatus` for Level-Zero backend. + ## Terms - Device-side Runtime - part of device-code, which is supplied by Device-side Compiler. - - Low-level Runtime - the backend/runtime, behind DPCPP Runtime. - - Device-side Compiler - compiler which generates device-native bitcode based - on input SPIR-V image. + - Device-side Compiler - compiler which generates device-native binary image + based on input SPIR-V image. + - Low-level Runtime - the backend/runtime behind DPCPP Runtime. - Accessor metadata - parts of accessor representation at device-side: pointer, ranges, offset. + ## How it works? For the time being, `assert(expr)` macro ends up in call to `__devicelib_assert_fail` function. This function is part of [Device library extension](doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst#cl_intel_devicelib_cassert). Device code already contains call to the function. Currently, a device-binary is always linked against fallback implementation. + + +### Device-specific approach + Device-side compiler/linker provides their implementation of `__devicelib_assert_fail` and prefer this implementation over fallback one. @@ -119,9 +124,12 @@ at synchronization points. Refer to [OpenCL](doc/extensions/Assert/opencl.md) and [Level-Zero](doc/extensions/Assert/level-zero.md) extensions. + +### Device-agnostic approach + If Device-side Runtime doesn't support `__devicelib_assert_fail` then a buffer based approach comes in place. The approach doesn't require any support from -Device-side Runtime. Neither it does from Low-level Runtime. +Device-side Runtime and Compiler. Neither it does from Low-level Runtime. Within this approach, a dedicated assert buffer is allocated and implicit kernel argument is introduced. The argument is an accessor with `discard_read_write` @@ -132,7 +140,107 @@ accessor metadata from program scope variable and writes assert information to the assert buffer. Atomic operations are used in order to not overwrite existing information. -Storing and restoring of accessor metadata to/from program scope variable is -performed with help of builtins. Implementations of these builtins are -substituted by frontend. +Both storing of accessor metadata and writing assert failure is performed with +help of built-ins. Implementations of these builtins are substituted by +frontend. + +#### Built-ins operation + +Accessor is a pointer augmented with offset and two ranges (access range and +memory range). + +There are two built-ins provided by frontend: + * `__store_acc()` - to store accessor metadata into program-scope variable. + * `__store_assert_failure()` - to store flag about assert failure in a buffer + using the metadata stored in program-scope variable. + +The accessor should be stored to program scope variable in global address space +using atomic operations. Motivation for using atomic operations: the program may +contain several kernels and some of them could be running simultaneously on a +single device. + +The `__store_assert_failure()` built-in atomically sets a flag in a buffer. The +buffer is accessed using accessor metadata from program-scope variable. This +built-in return a boolean value which is `true` if the flag is set by this call +to `__store_assert_failure()` and `false` if the flag was already set. +Motivation for using atomic operation is the same as with `__store_acc()` +builtin. + +The following pseudo-code snippets shows how these built-ins are used. +First of all, assume the following code as user's one: +``` +void user_func(int X) { + assert(X && “X is nil”); +} + +int main() { + queue Q(...); + Q.submit([&] (handler& CGH) { + CGH.single_task([=] () { + do_smth(); + user_func(0); + do_smth_else(); + }); + }); + ... +} +``` + +The following LLVM IR pseudo code will be generated for the user's code: +``` +@AssertBufferPtr = global void* null +@AssertBufferAccessRange = ... +@AssertBufferMemoryRange = ... +@AssertBufferOffset = ... + +/// user's code +void user_func(int X) { +if (!(X && “X is nil")) { + __assert_fail(...); + } +} + +users_kernel(...) { + do_smth() + user_func(0); + do_smth_else(); +} + +/// a wrapped user's kernel +kernel(AssertBufferAccessor, OtherArguments...) { + __store_acc(AssertBufferAccessor); + users_kernel(OtherArguments...); +} + +/// __assert_fail belongs to Linux version of devicelib +void __assert_fail(...) { + ... + __devicelib_assert_fail(...); +} + +void __devicelib_assert_fail(Expr, File, Line, GlobalID, LocalID) { + ... + if (__store_assert_info()) + printf("Assertion `%s' failed in %s at line %i. GlobalID: %i, LocalID: %i", + Expr, File, Line, GlobalID, LocalID); +} + +/// The following are built-ins provided by frontend +void __store_acc(accessor) { + %1 = accessor.getPtr(); + store void * %1, void * @AssertBufferPtr +} + +bool __store_assert_info(...) { + AssertBAcc = __fetch_acc(); + // fill in data in AsBAcc + volatile int *Ptr = (volatile int *)AssertBAcc.getPtr(); + bool Expected = false; + bool Desired = true; + + return atomic_cas(Ptr, Expected, Desired, SequentialConsistentMemoryOrder); + // or it could be: + // return !atomic_exchange(Ptr, Desired, SequentialConsistentMemoryOrder); +} +``` From 1f8d9a91069806e79c63c8cb96143bb4b7195071 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Mon, 5 Apr 2021 18:00:58 +0300 Subject: [PATCH 05/45] Fix links Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index 956fd64732cf5..5c4deac3118aa 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -104,7 +104,7 @@ should be performed via calls to `clGetEventInfo` for OpenCL backend and ## How it works? For the time being, `assert(expr)` macro ends up in call to -`__devicelib_assert_fail` function. This function is part of [Device library extension](doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst#cl_intel_devicelib_cassert). +`__devicelib_assert_fail` function. This function is part of [Device library extension](extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst#cl_intel_devicelib_cassert). Device code already contains call to the function. Currently, a device-binary is always linked against fallback implementation. @@ -121,7 +121,7 @@ is responsible for: When detected, Low-level Runtime reports assert failure to DPCPP Runtime at synchronization points. -Refer to [OpenCL](doc/extensions/Assert/opencl.md) and [Level-Zero](doc/extensions/Assert/level-zero.md) +Refer to [OpenCL](extensions/Assert/opencl.md) and [Level-Zero](extensions/Assert/level-zero.md) extensions. From 2ee590c861704a559bfb2d1eb7a243b350b48d5e Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Mon, 5 Apr 2021 18:08:23 +0300 Subject: [PATCH 06/45] Clarify that assertion failure message is printed by DPCPP Runtime Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index 5c4deac3118aa..a40ef6cf9af3b 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -50,7 +50,8 @@ In this use-case every work-item with even X dimension will trigger assertion failure. Assertion failure should be reported via asynchronous exceptions. If asynchronous exception handler is set the failure is reported with `cl::sycl::event_error` exception. Otherwise, SYCL Runtime should trigger abort. -At least one failed assertion should be reported. +At least one failed assertion should be reported. The assertion failure message +is printed to `stderr` by SYCL Runtime. When multiple kernels are enqueued and both fail at assertion at least single assertion should be reported. From 77699a24713455836cdffec6a1feab30f4ba7ee3 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Tue, 6 Apr 2021 17:31:16 +0300 Subject: [PATCH 07/45] Clarify that fallback assert impl is synchronous Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 3 +++ 1 file changed, 3 insertions(+) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index a40ef6cf9af3b..ca85893ae4466 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -141,6 +141,9 @@ accessor metadata from program scope variable and writes assert information to the assert buffer. Atomic operations are used in order to not overwrite existing information. +DPCPP Runtime checks contents of the assert buffer for assert failure flag after +kernel finishes. + Both storing of accessor metadata and writing assert failure is performed with help of built-ins. Implementations of these builtins are substituted by frontend. From 001a5736e4c3f7e269b889585d1b401879e681a6 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Tue, 6 Apr 2021 18:22:30 +0300 Subject: [PATCH 08/45] Fix typo in level-zero ext draft Signed-off-by: Sergey Kanaev --- sycl/doc/extensions/Assert/level-zero.md | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/sycl/doc/extensions/Assert/level-zero.md b/sycl/doc/extensions/Assert/level-zero.md index 265ff5858d83f..f3b8e402db052 100644 --- a/sycl/doc/extensions/Assert/level-zero.md +++ b/sycl/doc/extensions/Assert/level-zero.md @@ -4,8 +4,9 @@ This extension enables detection of assert failure of kernel. # New enum value -`ze_result_t` enumeration should be augmented with `ZE_RESULT_ABORTED` enum -element. This enum value indicated a detected assert failure at device-side. +`ze_result_t` enumeration should be augmented with `ZE_RESULT_ASSERT_FAILED` +enum element. This enum value indicated a detected assert failure at +device-side. # Changed API @@ -15,5 +16,5 @@ ze_result Result = zeEventQueryStatus(Event); ``` If kernel failed an assertion `zeEventQueryStatus` should return -`ZE_RESULT_ABORTED`. +`ZE_RESULT_ASSERT_FAILED`. From 32b647987d3db0072f1bac298b726c8a867467ec Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 7 Apr 2021 18:29:40 +0300 Subject: [PATCH 09/45] Address some review comments. Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 85 ++++++++++++++++++++++++++++------------------ 1 file changed, 52 insertions(+), 33 deletions(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index ca85893ae4466..d57c55658188e 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -2,19 +2,20 @@ **IMPORTANT**: This document is a draft. -During debugging of kernel code user may put assertions here and there. -The expected behaviour of assertion failure at host is application abort. -Our choice for device-side assertions is asynchronous exception in order to -allow for extensibility. - -The user is free to disable assertions by defining `NDEBUG` macro at -compile-time. +Using the standard C++ `assert` API ("assertions") is an important debugging +technique widely used by developers. This document describes the design of +supporting assertions within SYCL device code. +The basic approach we chose is delivering device-side assertions as host-side +asynchronous exceptions, which allows further extensibility, such as better +error handling or potential recovery. +As usual, device-side assertions can be disabled by defining `NDEBUG` macro at +compile time. ## Use-case example ``` -using namespace cl::sycl; +using namespace sycl; auto ErrorHandler = [] (exception_list Exs) { for (exception_ptr const& E : Exs) { try { @@ -49,12 +50,13 @@ int main() { In this use-case every work-item with even X dimension will trigger assertion failure. Assertion failure should be reported via asynchronous exceptions. If asynchronous exception handler is set the failure is reported with -`cl::sycl::event_error` exception. Otherwise, SYCL Runtime should trigger abort. -At least one failed assertion should be reported. The assertion failure message -is printed to `stderr` by SYCL Runtime. +`sycl::event_error` exception. Otherwise, SYCL Runtime should trigger abort. +Even though multiple failures of the same or different assertions can happen in +multiple workitems, implementation is required to deliver only one. The +assertion failure message is printed to `stderr` by SYCL Runtime. -When multiple kernels are enqueued and both fail at assertion at least single -assertion should be reported. +When multiple kernels are enqueued and more than one fail at assertion, at least +single assertion should be reported. ## User requirements @@ -68,10 +70,14 @@ From user's point of view there are the following requirements: | 3 | Stop under debugger | When debugger is attached, break at assertion point | Highly desired | | 4 | Reliability | Assert failure should be reported regardless of kernel deadlock | Highly desired | +Implementations without enough capabilities to implement fourth requirement are +allowed to realize the fallback approach described below, which does not +guarantee assertion failure delivery to host, but is still useful in many +practical cases. -## Contents of `cl::sycl::event_error` +## Contents of `sycl::event_error` -Interface of `cl::sycl::event_error` should look like: +Interface of `sycl::event_error` should look like: ``` class event_error : public runtime_error { public: @@ -87,50 +93,63 @@ public: Regardless of whether asynchronous exception handler is set or not, there's an action to be performed by SYCL Runtime. To achieve this, information about assert failure should be propagated from device-side to SYCL Runtime. This -should be performed via calls to `clGetEventInfo` for OpenCL backend and -`zeEventQueryStatus` for Level-Zero backend. +should be performed via calls to `piEventGetInfo`. This Plugin Interface call +"lowers" to `clGetEventInfo` for OpenCL backend and `zeEventQueryStatus` for +Level-Zero backend. ## Terms - - Device-side Runtime - part of device-code, which is supplied by Device-side - Compiler. - - Device-side Compiler - compiler which generates device-native binary image + - Device-side Runtime - runtime library supplied by the Native Device Compiler + and running on the device. + - Native Device Compiler - compiler which generates device-native binary image based on input SPIR-V image. - - Low-level Runtime - the backend/runtime behind DPCPP Runtime. + - Low-level Runtime - the backend/runtime behind DPCPP Runtime attached via the + Plugin Interface. - Accessor metadata - parts of accessor representation at device-side: pointer, ranges, offset. ## How it works? -For the time being, `assert(expr)` macro ends up in call to -`__devicelib_assert_fail` function. This function is part of [Device library extension](extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst#cl_intel_devicelib_cassert). -Device code already contains call to the function. Currently, a device-binary -is always linked against fallback implementation. +`assert(expr)` macro ends up in call to `__devicelib_assert_fail`. This function +is part of [Device library extension](extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst#cl_intel_devicelib_cassert). + +Implementation of this function is supplied by Native Device Compiler for +safe approach or by DPCPP Compiler for fallback one. +Due to lack of support of online linking in Level-Zero, the application is +linked against fallback implementation of `__devicelib_assert_fail`. Hence, +Native Device Compilers should prefer their implementation instead of the one +provided in incoming SPIR-V/LLVM IR binary. -### Device-specific approach -Device-side compiler/linker provides their implementation of `__devicelib_assert_fail` -and prefer this implementation over fallback one. +### Safe approach -If Device-side Runtime supports `__devicelib_assert_fail` then Low-Level Runtime -is responsible for: +This is the preferred approach and implementations should use it when possible. +It guarantees assertion failure notification delivery to the host regardless of +kernel behavior which hit the assertion. + +The Native Device Compiler is responsible for providing implementation of +`__devicelib_assert_fail` which completely hides details of communication +between the device code and the Low-Level Runtime from the SYCL device compiler +and runtime. The Low-Level Runtime is responsible for: - detecting if assert failure took place; - flushing assert message to `stderr` on host. + When detected, Low-level Runtime reports assert failure to DPCPP Runtime -at synchronization points. +via events objects. Refer to [OpenCL](extensions/Assert/opencl.md) and [Level-Zero](extensions/Assert/level-zero.md) extensions. -### Device-agnostic approach +### Fallback approach If Device-side Runtime doesn't support `__devicelib_assert_fail` then a buffer based approach comes in place. The approach doesn't require any support from -Device-side Runtime and Compiler. Neither it does from Low-level Runtime. +Device-side Runtime and Native Device Compiler. Neither it does from Low-level +Runtime. Within this approach, a dedicated assert buffer is allocated and implicit kernel argument is introduced. The argument is an accessor with `discard_read_write` From b8637c2003cd558578886c3de821a6125eb6fe65 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Thu, 8 Apr 2021 16:30:00 +0300 Subject: [PATCH 10/45] Add exception extension Signed-off-by: Sergey Kanaev --- .../SYCL_INTEL_assert_exception.asciidoc | 109 ++++++++++++++++++ 1 file changed, 109 insertions(+) create mode 100644 sycl/doc/extensions/Assert/SYCL_INTEL_assert_exception.asciidoc diff --git a/sycl/doc/extensions/Assert/SYCL_INTEL_assert_exception.asciidoc b/sycl/doc/extensions/Assert/SYCL_INTEL_assert_exception.asciidoc new file mode 100644 index 0000000000000..691548bfa9502 --- /dev/null +++ b/sycl/doc/extensions/Assert/SYCL_INTEL_assert_exception.asciidoc @@ -0,0 +1,109 @@ += SYCL_INTEL_assert_exception + +:source-highlighter: coderay +:coderay-linenums-mode: table + +// This section needs to be after the document title. +:doctype: book +:toc2: +:toc: left +:encoding: utf-8 +:lang: en + +:blank: pass:[ +] + +// Set the default source code type in this document to C++, +// for syntax highlighting purposes. This is needed because +// docbook uses c++ and html5 uses cpp. +:language: {basebackend@docbook:c++:cpp} + +// This is necessary for asciidoc, but not for asciidoctor +:cpp: C++ + +== Introduction +IMPORTANT: This specification is a draft. + +NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are +trademarks of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. +used by permission by Khronos. + +NOTE: This document is better viewed when rendered as html with asciidoctor. +GitHub does not render image icons. + +This document describes an extension to rename device-specific kernel queries +to better describe the operations performed. + +== Name Strings + ++SYCL_INTEL_assert_exception+ + +== Notice + +Copyright (c) 2021 Intel Corporation. All rights reserved. + +== Status + +Working Draft + +This is a preview extension specification, intended to provide early access to +a feature for review and community feedback. When the feature matures, this +specification may be released as a formal extension. + +Because the interfaces defined by this specification are not final and are +subject to change they are not intended to be used by shipping software +products. + +== Version + +Built On: {docdate} + +Revision: 1 + +== Dependencies + +This extension is written against the SYCL 2020 specification, Revision 3. + +== Overview + +== Modifications of SYCL 2020 Specification + +=== Change Section 4.13.2 Exception class interface + +Add enum member `assert` to the `errc` enum class: + +[source,c++,`sycl::kernel`,linenums] +---- +assert = /* implementation defined */ +---- + +==== Change table 136 Values of `errc` enum + +Add row `assert`: + +[width="40%",frame="topbot",options="header,footer"] +|====================== +|Standard SYCL Error Codes |Description +|`assert` | Assert failure had happened in device code during kernel execution +|====================== + +== Issues + +None. + +== Revision History + +[cols="5,15,15,70"] +[grid="rows"] +[options="header"] +|======================================== +|Rev|Date|Author|Changes +|1|2021-04-08|Sergey Kanaev|*Initial public working draft* +|======================================== + +//************************************************************************ +//Other formatting suggestions: +// +//* Use *bold* text for host APIs, or [source] syntax highlighting. +//* Use +mono+ text for device APIs, or [source] syntax highlighting. +//* Use +mono+ text for extension names, types, or enum values. +//* Use _italics_ for parameters. +//************************************************************************ From b0cd85ff063a3bb5f85d38aa119f694600f3c476 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Thu, 8 Apr 2021 17:20:59 +0300 Subject: [PATCH 11/45] Use error-code instead of distinct exception. Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 36 ++++++++---------------------------- 1 file changed, 8 insertions(+), 28 deletions(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index d57c55658188e..058fde5126a76 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -48,12 +48,11 @@ int main() { ``` In this use-case every work-item with even X dimension will trigger assertion -failure. Assertion failure should be reported via asynchronous exceptions. If -asynchronous exception handler is set the failure is reported with -`sycl::event_error` exception. Otherwise, SYCL Runtime should trigger abort. +failure. Assertion failure should be reported via asynchronous exceptions with +[`assert` error code](extensions/Assert/SYCL_INTEL_assert_exception.asciidoc). Even though multiple failures of the same or different assertions can happen in multiple workitems, implementation is required to deliver only one. The -assertion failure message is printed to `stderr` by SYCL Runtime. +assertion failure message is printed to `stderr` by DPCPP Runtime. When multiple kernels are enqueued and more than one fail at assertion, at least single assertion should be reported. @@ -75,28 +74,6 @@ allowed to realize the fallback approach described below, which does not guarantee assertion failure delivery to host, but is still useful in many practical cases. -## Contents of `sycl::event_error` - -Interface of `sycl::event_error` should look like: -``` -class event_error : public runtime_error { -public: - event_error() = default; - - event_error(const char *Msg, cl_int Err) - : event_error(string_class(Msg), Err) {} - - event_error(const string_class &Msg, cl_int Err) : runtime_error(Msg, Err) {} -}; -``` - -Regardless of whether asynchronous exception handler is set or not, there's an -action to be performed by SYCL Runtime. To achieve this, information about -assert failure should be propagated from device-side to SYCL Runtime. This -should be performed via calls to `piEventGetInfo`. This Plugin Interface call -"lowers" to `clGetEventInfo` for OpenCL backend and `zeEventQueryStatus` for -Level-Zero backend. - ## Terms @@ -118,7 +95,7 @@ is part of [Device library extension](extensions/C-CXX-StandardLibrary/DeviceLib Implementation of this function is supplied by Native Device Compiler for safe approach or by DPCPP Compiler for fallback one. -Due to lack of support of online linking in Level-Zero, the application is +NB: Due to lack of support of online linking in Level-Zero, the application is linked against fallback implementation of `__devicelib_assert_fail`. Hence, Native Device Compilers should prefer their implementation instead of the one provided in incoming SPIR-V/LLVM IR binary. @@ -138,7 +115,10 @@ and runtime. The Low-Level Runtime is responsible for: - flushing assert message to `stderr` on host. When detected, Low-level Runtime reports assert failure to DPCPP Runtime -via events objects. +via events objects. To achieve this, information about assert failure should be +propagated from device-side to SYCL Runtime. This should be performed via calls +to `piEventGetInfo`. This Plugin Interface call "lowers" to `clGetEventInfo` for +OpenCL backend and `zeEventQueryStatus` for Level-Zero backend. Refer to [OpenCL](extensions/Assert/opencl.md) and [Level-Zero](extensions/Assert/level-zero.md) extensions. From 8c036486b0864399d1b59144d03f390cd7972920 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Fri, 9 Apr 2021 17:26:46 +0300 Subject: [PATCH 12/45] [SYCL] Add OpenCL extension for assert error code Signed-off-by: Sergey Kanaev --- .../cl_intel_assert_return_code.asciidoc | 99 +++++++++++++++++++ 1 file changed, 99 insertions(+) create mode 100644 sycl/doc/extensions/Assert/cl_intel_assert_return_code.asciidoc diff --git a/sycl/doc/extensions/Assert/cl_intel_assert_return_code.asciidoc b/sycl/doc/extensions/Assert/cl_intel_assert_return_code.asciidoc new file mode 100644 index 0000000000000..b7eec45d0a26f --- /dev/null +++ b/sycl/doc/extensions/Assert/cl_intel_assert_return_code.asciidoc @@ -0,0 +1,99 @@ +cl_intel_assert_return_code +====================================== + +// This section needs to be after the document title. +:doctype: book +:toc2: +:toc: left +:encoding: utf-8 +:lang: en + +:blank: pass:[ +] + +// Set the default source code type in this document to C++, +// for syntax highlighting purposes. This is needed because +// docbook uses c++ and html5 uses cpp. +:language: {basebackend@docbook:c++:cpp} + +== Name Strings + ++cl_intel_assert_return_code+ + +== Notice + +Copyright (c) 2021 Intel Corporation. All rights reserved. + +== Status + +Working Draft + +This is a preview extension specification, intended to provide early access to +a feature for review and community feedback. When the feature matures, this +specification may be released as a formal extension. + +Because the interfaces defined by this specification are not final and are +subject to change they are not intended to be used by shipping software +products. + +== Version + +Built On: {docdate} + +Revision: 1 + +== Dependencies + +This extension is written against the OpenCL Specification Version 1.0, Revision 48. + +This extension requires OpenCL 1.0 or later. + +== Overview + +This extension allows OpenCL 1.x and 2.x devices to notify host that assert had +happened. + +== New error code + +[source] +---- +CL_ASSERT_FAILURE +---- + +Negative value of this error code should be set into `param_value` of +`clGetEventInfo` as described in table 5.15 "clGetEventInfo prameter queries" if +assert failure took place in device-code during kernel execution. + +An example: +[source] +---- +cl_event Event; // describes an event of kernel been submitted previously +cl_int Result; +size_t ResultSize; + +clGetEventInfo(Event, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(Result), &Result, &ResultSize); +---- + +If kernel failed an assertion `clGetEventInfo` should put `CL_ASSERT_FAILURE` in +`Result`. + +== Issues + +None. + +== Revision History + +[cols="5,15,15,70"] +[grid="rows"] +[options="header"] +|======================================== +|Rev|Date|Author|Changes +|1|2021-04-09|Sergey Kanaev|*Initial public working draft* +|======================================== + +//************************************************************************ +//Other formatting suggestions: +// +//* Use *bold* text for host APIs, or [source] syntax highlighting. +//* Use +mono+ text for device APIs, or [source] syntax highlighting. +//* Use +mono+ text for extension names, types, or enum values. +//* Use _italics_ for parameters. +//************************************************************************ From 121c945bbbe2cff52d4766127de3bbddb9df4d68 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Fri, 9 Apr 2021 17:50:17 +0300 Subject: [PATCH 13/45] [SYCL] Add Level-Zero extension for assert error code Signed-off-by: Sergey Kanaev --- .../ze_intel_assert_return_code.asciidoc | 124 ++++++++++++++++++ 1 file changed, 124 insertions(+) create mode 100644 sycl/doc/extensions/Assert/ze_intel_assert_return_code.asciidoc diff --git a/sycl/doc/extensions/Assert/ze_intel_assert_return_code.asciidoc b/sycl/doc/extensions/Assert/ze_intel_assert_return_code.asciidoc new file mode 100644 index 0000000000000..b56937f9ba0d3 --- /dev/null +++ b/sycl/doc/extensions/Assert/ze_intel_assert_return_code.asciidoc @@ -0,0 +1,124 @@ +ze_intel_assert_return_code +====================================== + +// This section needs to be after the document title. +:doctype: book +:toc2: +:toc: left +:encoding: utf-8 +:lang: en + +:blank: pass:[ +] + +// Set the default source code type in this document to C++, +// for syntax highlighting purposes. This is needed because +// docbook uses c++ and html5 uses cpp. +:language: {basebackend@docbook:c++:cpp} + +== Name Strings + ++ze_intel_assert_return_code+ + +== Notice + +Copyright (c) 2021 Intel Corporation. All rights reserved. + +== Status + +Working Draft + +This is a preview extension specification, intended to provide early access to +a feature for review and community feedback. When the feature matures, this +specification may be released as a formal extension. + +Because the interfaces defined by this specification are not final and are +subject to change they are not intended to be used by shipping software +products. + +== Version + +Built On: {docdate} + +Revision: 1 + +== Dependencies + +This extension is written against the Level-Zero Specification Version 1.1.2. + +== Overview + +This extension allows Level-Zero devices to notify host that assert had +happened. + +== New enumeration value + +`ze_result_t`: + +[source] +---- +ZE_RESULT_ASSERT_FAILED +---- + +This value should be returned by `zeEventQueryStatus` if assert failure took +place in device-code during kernel execution. + +An example: +[source] +---- +ze_event_handle_t Event; // describes an event of kernel been submitted previously +ze_result Result = zeEventQueryStatus(Event); +---- + +If kernel failed an assertion `zeEventQueryStatus` should return +`ZE_RESULT_ASSERT_FAILED`. + + +== Modifications to Level-Zero API + +(Add to Section API Documentation / Core API / Common / Common Enums / `ze_result_t`) :: ++ +-- +`ZE_RESULT_ASSERT_FAILED = 0x70000006` + +[Core] Assert failure took place in device-code during kernel execution. +-- + +(Add to section API Documentation / Core API / Event / Event Functions / `zeEventQueryStatus`) :: ++ +-- +Return: + +`ZE_RESULT_ASSERT_FAILED` +-- + +An example: +[source] +---- +cl_event Event; // describes an event of kernel been submitted previously +cl_int Result; +size_t ResultSize; + +clGetEventInfo(Event, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(Result), &Result, &ResultSize); +---- + +If kernel failed an assertion `clGetEventInfo` should put `CL_ASSERT_FAILURE` in +`Result`. + +== Issues + +None. + +== Revision History + +[cols="5,15,15,70"] +[grid="rows"] +[options="header"] +|======================================== +|Rev|Date|Author|Changes +|1|2021-04-09|Sergey Kanaev|*Initial public working draft* +|======================================== + +//************************************************************************ +//Other formatting suggestions: +// +//* Use *bold* text for host APIs, or [source] syntax highlighting. +//* Use +mono+ text for device APIs, or [source] syntax highlighting. +//* Use +mono+ text for extension names, types, or enum values. +//* Use _italics_ for parameters. +//************************************************************************ From a4b48849bbbbc290f3cb0168894a250579fe1bc0 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Fri, 9 Apr 2021 17:52:28 +0300 Subject: [PATCH 14/45] Remove draft files Signed-off-by: Sergey Kanaev --- sycl/doc/extensions/Assert/level-zero.md | 20 -------------------- sycl/doc/extensions/Assert/opencl.md | 22 ---------------------- 2 files changed, 42 deletions(-) delete mode 100644 sycl/doc/extensions/Assert/level-zero.md delete mode 100644 sycl/doc/extensions/Assert/opencl.md diff --git a/sycl/doc/extensions/Assert/level-zero.md b/sycl/doc/extensions/Assert/level-zero.md deleted file mode 100644 index f3b8e402db052..0000000000000 --- a/sycl/doc/extensions/Assert/level-zero.md +++ /dev/null @@ -1,20 +0,0 @@ -# Overview - -This extension enables detection of assert failure of kernel. - -# New enum value - -`ze_result_t` enumeration should be augmented with `ZE_RESULT_ASSERT_FAILED` -enum element. This enum value indicated a detected assert failure at -device-side. - -# Changed API - -``` -ze_event_handle_t Event; // describes an event of kernel been submitted previously -ze_result Result = zeEventQueryStatus(Event); -``` - -If kernel failed an assertion `zeEventQueryStatus` should return -`ZE_RESULT_ASSERT_FAILED`. - diff --git a/sycl/doc/extensions/Assert/opencl.md b/sycl/doc/extensions/Assert/opencl.md deleted file mode 100644 index 50ad0b7db0897..0000000000000 --- a/sycl/doc/extensions/Assert/opencl.md +++ /dev/null @@ -1,22 +0,0 @@ -# Overview - -This extension enables detection of assert failure of kernel. - -# New error code - -`CL_ASSERT_FAILURE` is added to indicate a detected assert failure at -device-side. - -# Changed API - -``` -cl_event Event; // describes an event of kernel been submitted previously -cl_int Result; -size_t ResultSize; - -clGetEventInfo(Event, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(Result), &Result, &ResultSize); -``` - -If kernel failed an assertion `clGetEventInfo` should put `CL_ASSERT_FAILURE` -in `Result`. - From c06db5f013f337c5ef76d56d5ade53e7287d273e Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Fri, 9 Apr 2021 17:55:16 +0300 Subject: [PATCH 15/45] Remove unwanted part Signed-off-by: Sergey Kanaev --- .../Assert/ze_intel_assert_return_code.asciidoc | 13 ------------- 1 file changed, 13 deletions(-) diff --git a/sycl/doc/extensions/Assert/ze_intel_assert_return_code.asciidoc b/sycl/doc/extensions/Assert/ze_intel_assert_return_code.asciidoc index b56937f9ba0d3..85db891ad5bcf 100644 --- a/sycl/doc/extensions/Assert/ze_intel_assert_return_code.asciidoc +++ b/sycl/doc/extensions/Assert/ze_intel_assert_return_code.asciidoc @@ -87,19 +87,6 @@ Return: + `ZE_RESULT_ASSERT_FAILED` -- -An example: -[source] ----- -cl_event Event; // describes an event of kernel been submitted previously -cl_int Result; -size_t ResultSize; - -clGetEventInfo(Event, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(Result), &Result, &ResultSize); ----- - -If kernel failed an assertion `clGetEventInfo` should put `CL_ASSERT_FAILURE` in -`Result`. - == Issues None. From a99368bfb0cd82142852296370d4181560897d6f Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Fri, 9 Apr 2021 18:00:19 +0300 Subject: [PATCH 16/45] Add limitations on submit to same queue after exception thrown. Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index 058fde5126a76..a677b0113aee3 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -100,6 +100,13 @@ linked against fallback implementation of `__devicelib_assert_fail`. Hence, Native Device Compilers should prefer their implementation instead of the one provided in incoming SPIR-V/LLVM IR binary. +Limitations for user: + - DPCPP RT, Low-level RT and device state is unknown after throwing of "assert" + asynchronous exception; + - "assert" asynchronous exception might not be recoverable; + - there might not be guarantees on enqueueing commands (kernel, copy, etc.) to + same queue/context: guarantees may vary with device/Low-level RT. + ### Safe approach From 78d7fcbd5ef4765078eaeeaab9a723ec68a7ea7b Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Fri, 9 Apr 2021 18:04:21 +0300 Subject: [PATCH 17/45] Add format of assert message Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index a677b0113aee3..ce415ff9eef83 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -92,6 +92,11 @@ practical cases. `assert(expr)` macro ends up in call to `__devicelib_assert_fail`. This function is part of [Device library extension](extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst#cl_intel_devicelib_cassert). +Format of assert failure message, printed to `stderr` is the following: +``` +:: : global id: [,,], local id: [,,] Assertion `` failed. +``` + Implementation of this function is supplied by Native Device Compiler for safe approach or by DPCPP Compiler for fallback one. From 6882e95a17823c92e29925d0e70a84b3f47b7ac2 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Fri, 9 Apr 2021 18:18:36 +0300 Subject: [PATCH 18/45] Clarify where kernel wrapping takes place Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 35 +++++++++++++++++++++++++++++++++++ 1 file changed, 35 insertions(+) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index ce415ff9eef83..2789e200c37b3 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -159,6 +159,41 @@ Both storing of accessor metadata and writing assert failure is performed with help of built-ins. Implementations of these builtins are substituted by frontend. +User's kernel is executed through a wrapper. Wrapping takes place in DPCPP +Runtime headers in a following manner: + +``` +class handler { + +template parallel_for(KernelFunc, Range) { +#ifndef NDEBUG + // Assert required +  if (!MQueue->get_device()->assert_fail_supported()) { +    using KName2 = class ASSERT_WRAPPER_NAME(KernelName); +     +    auto AssertBufferAcc = MQueue->get_context()->getAssertBufferAccessor(this); + +    parallel_for_impl( +      Range, +      [=](Item) { +        __store_acc(AssertBuffAcc); +        KernelFunc(Item); +      }); +  } else { +#endif + + // (No assert required) OR (Assert supported by device) +     // ordinary enqueue process + +#ifndef NDEBUG + } +#endif +} + +} +``` + + #### Built-ins operation Accessor is a pointer augmented with offset and two ranges (access range and From 32663e03410db9685c2eb6841621d558a9267e3f Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Tue, 13 Apr 2021 18:36:35 +0300 Subject: [PATCH 19/45] Changes to SYCL specification Signed-off-by: Sergey Kanaev --- .../SYCL_INTEL_assert_exception.asciidoc | 109 --------------- .../Assert/SYCL_ONEAPI_ASSERT.asciidoc | 131 ++++++++++++++++++ .../DeviceLibExtensions.rst | 2 + 3 files changed, 133 insertions(+), 109 deletions(-) delete mode 100644 sycl/doc/extensions/Assert/SYCL_INTEL_assert_exception.asciidoc create mode 100644 sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc diff --git a/sycl/doc/extensions/Assert/SYCL_INTEL_assert_exception.asciidoc b/sycl/doc/extensions/Assert/SYCL_INTEL_assert_exception.asciidoc deleted file mode 100644 index 691548bfa9502..0000000000000 --- a/sycl/doc/extensions/Assert/SYCL_INTEL_assert_exception.asciidoc +++ /dev/null @@ -1,109 +0,0 @@ -= SYCL_INTEL_assert_exception - -:source-highlighter: coderay -:coderay-linenums-mode: table - -// This section needs to be after the document title. -:doctype: book -:toc2: -:toc: left -:encoding: utf-8 -:lang: en - -:blank: pass:[ +] - -// Set the default source code type in this document to C++, -// for syntax highlighting purposes. This is needed because -// docbook uses c++ and html5 uses cpp. -:language: {basebackend@docbook:c++:cpp} - -// This is necessary for asciidoc, but not for asciidoctor -:cpp: C++ - -== Introduction -IMPORTANT: This specification is a draft. - -NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are -trademarks of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. -used by permission by Khronos. - -NOTE: This document is better viewed when rendered as html with asciidoctor. -GitHub does not render image icons. - -This document describes an extension to rename device-specific kernel queries -to better describe the operations performed. - -== Name Strings - -+SYCL_INTEL_assert_exception+ - -== Notice - -Copyright (c) 2021 Intel Corporation. All rights reserved. - -== Status - -Working Draft - -This is a preview extension specification, intended to provide early access to -a feature for review and community feedback. When the feature matures, this -specification may be released as a formal extension. - -Because the interfaces defined by this specification are not final and are -subject to change they are not intended to be used by shipping software -products. - -== Version - -Built On: {docdate} + -Revision: 1 - -== Dependencies - -This extension is written against the SYCL 2020 specification, Revision 3. - -== Overview - -== Modifications of SYCL 2020 Specification - -=== Change Section 4.13.2 Exception class interface - -Add enum member `assert` to the `errc` enum class: - -[source,c++,`sycl::kernel`,linenums] ----- -assert = /* implementation defined */ ----- - -==== Change table 136 Values of `errc` enum - -Add row `assert`: - -[width="40%",frame="topbot",options="header,footer"] -|====================== -|Standard SYCL Error Codes |Description -|`assert` | Assert failure had happened in device code during kernel execution -|====================== - -== Issues - -None. - -== Revision History - -[cols="5,15,15,70"] -[grid="rows"] -[options="header"] -|======================================== -|Rev|Date|Author|Changes -|1|2021-04-08|Sergey Kanaev|*Initial public working draft* -|======================================== - -//************************************************************************ -//Other formatting suggestions: -// -//* Use *bold* text for host APIs, or [source] syntax highlighting. -//* Use +mono+ text for device APIs, or [source] syntax highlighting. -//* Use +mono+ text for extension names, types, or enum values. -//* Use _italics_ for parameters. -//************************************************************************ diff --git a/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc b/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc new file mode 100644 index 0000000000000..8b79bf86fa967 --- /dev/null +++ b/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc @@ -0,0 +1,131 @@ += SYCL_EXT_ONEAPI_ASSERT + +:source-highlighter: coderay +:coderay-linenums-mode: table + +// This section needs to be after the document title. +:doctype: book +:toc2: +:toc: left +:encoding: utf-8 +:lang: en + +:blank: pass:[ +] + +// Set the default source code type in this document to C++, +// for syntax highlighting purposes. This is needed because +// docbook uses c++ and html5 uses cpp. +:language: {basebackend@docbook:c++:cpp} + +// This is necessary for asciidoc, but not for asciidoctor +:cpp: C++ + +== Notice +IMPORTANT: This specification is a draft. + +Copyright (c) 2021-2021 Intel Corporation. All rights reserved. + +NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are +trademarks of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. +used by permission by Khronos. + +NOTE: This document is better viewed when rendered as html with asciidoctor. +GitHub does not render image icons. + +== Dependencies + +This extension is written against the SYCL 2020 specification, Revision 3. + +== Status + +Working Draft + +This is a preview extension specification, intended to provide early access to +a feature for review and community feedback. When the feature matures, this +specification may be released as a formal extension. + +Because the interfaces defined by this specification are not final and are +subject to change they are not intended to be used by shipping software +products. + +== Introduction +This extension adds the ability for device code to call the C++ `assert()` +macro. The behavior of `assert()` in device code is similar to its behavior in +host code. If the asserted condition is false, a message is printed to `stderr` +and then the program typically aborts. The mechanism for aborting the program is +different, though. Whereas the host version calls `std::abort()`, the device +version causes an asynchronous SYCL `exception` with the +`errc::ext_oneapi_assert` error code to be thrown. The application can catch +this exception like any other asynchronous exception that is thrown from an +executing kernel. The numeric value of this enumeration element is defined by +implementation. + + +Catching the asynchronous exception with `sycl::errc::ext_oneapi_assert` error +code means that assert failure had happened in device code during kernel +execution and the assert message is already printed to `stderr`. Format of the +assert message is the following: +``` +:: : global id: [,,], local id: [,,] Assertion `` failed. +``` + +It is unspecified whether a failing `assert()` returns to its caller before the +kernel terminates. If a failing call returns, the device code may need to +continue execution without deadlocking for the assertion message to be printed +or for the exception to be thrown. + +The contents of the exception's `e.what()` string is unspecified. Since the +assertion message is already printed to `stderr` by the time the exception is +thrown, the `e.what()` string may not include the location of the failing +assertion. + +The `assert()` macro is defined in system include headers, not in SYCL headers. +On most of systems it is `` and/or `` header files. +The user can can disable assertions in device code by defining the `NDEBUG` +preprocessor macro at the point when `` and `/` +are included. + + +== Feature test macro + +This extension provides a feature-test macro as described in the core SYCL +specification section 6.3.3 "Feature test macros". Therefore, an implementation +supporting this extension must predefine the macro `SYCL_EXT_ONEAPI_ASSERT` to +one of the values defined in the table below. Applications can test for the +existence of this macro to determine if the implementation supports this +feature, or applications can test the macro’s value to determine which of the +extension’s APIs the implementation supports. + +[%header,cols="1,5"] +|=== +|Value |Description +|1 |Initial extension version. Base features are supported. +|=== + +== Version + +Built On: {docdate} + +Revision: 1 + +== Issues + +None. + +== Revision History + +[cols="5,15,15,70"] +[grid="rows"] +[options="header"] +|======================================== +|Rev|Date|Author|Changes +|1|2021-04-08|Sergey Kanaev, Gregory M Lueck |*Initial public working draft* +|======================================== + +//************************************************************************ +//Other formatting suggestions: +// +//* Use *bold* text for host APIs, or [source] syntax highlighting. +//* Use +mono+ text for device APIs, or [source] syntax highlighting. +//* Use +mono+ text for extension names, types, or enum values. +//* Use _italics_ for parameters. +//************************************************************************ diff --git a/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst b/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst index 8b8b98d7a12bb..62ed492a76ac4 100644 --- a/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst +++ b/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst @@ -33,6 +33,8 @@ 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. +See also: [assert extension](../Assert/SYCL_ONEAPI_ASSERT.asciidoc) + cl_intel_devicelib_math ========================== From 2b84a83dce857993b7dcef396f0ff9a343ec1c9b Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Tue, 13 Apr 2021 18:58:59 +0300 Subject: [PATCH 20/45] Elaborate on limitations Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index 2789e200c37b3..22d626adcd1f4 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -105,12 +105,12 @@ linked against fallback implementation of `__devicelib_assert_fail`. Hence, Native Device Compilers should prefer their implementation instead of the one provided in incoming SPIR-V/LLVM IR binary. -Limitations for user: - - DPCPP RT, Low-level RT and device state is unknown after throwing of "assert" - asynchronous exception; - - "assert" asynchronous exception might not be recoverable; - - there might not be guarantees on enqueueing commands (kernel, copy, etc.) to - same queue/context: guarantees may vary with device/Low-level RT. +Limitations for user after catching the "assert" asynchronous exception: + - When using GPU device and the kernel hangs/crashes the subsequent enqueues + will fail; +When using CPU devices the user can proceed with enqueues to the same +device/queue/context. +DPCPP Runtime remains in valid state after "assert" exception been thrown. ### Safe approach From 423107b3f22ec85fdead940dce4b1b9525fea361 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 14 Apr 2021 10:19:58 +0300 Subject: [PATCH 21/45] Fix link Signed-off-by: Sergey Kanaev --- .../extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst b/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst index 62ed492a76ac4..1c370e57ad89c 100644 --- a/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst +++ b/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst @@ -33,7 +33,8 @@ 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. -See also: [assert extension](../Assert/SYCL_ONEAPI_ASSERT.asciidoc) +See also: assert_extension_. +.. _assert_extension: ../Assert/SYCL_ONEAPI_ASSERT.asciidoc) cl_intel_devicelib_math ========================== From 76115114280afcd1a6fec0afd69171c4656715cc Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 14 Apr 2021 14:46:58 +0300 Subject: [PATCH 22/45] Add sequence describing how DPCPP RT gets to know about assert failure Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 37 ++++++++++++++++++++++++++++++++++++- 1 file changed, 36 insertions(+), 1 deletion(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index 22d626adcd1f4..8a6767af12c15 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -108,12 +108,19 @@ provided in incoming SPIR-V/LLVM IR binary. Limitations for user after catching the "assert" asynchronous exception: - When using GPU device and the kernel hangs/crashes the subsequent enqueues will fail; + When using CPU devices the user can proceed with enqueues to the same device/queue/context. DPCPP Runtime remains in valid state after "assert" exception been thrown. -### Safe approach +### Current violation + +While throwing an asynchronous exception is quite an extensible way, for the +time being DPCPP Runtime merely calls `abort()`. + + +## Safe approach This is the preferred approach and implementations should use it when possible. It guarantees assertion failure notification delivery to the host regardless of @@ -135,6 +142,20 @@ OpenCL backend and `zeEventQueryStatus` for Level-Zero backend. Refer to [OpenCL](extensions/Assert/opencl.md) and [Level-Zero](extensions/Assert/level-zero.md) extensions. +The following sequence of events describes how user code gets notified: + - Device side: + 1. Assert fails in device-code in kernel + // It's not defined if GPU thread stops execution + // Other GPU threads are left untouched + 2. Specialized version of `__devicelib_assert_fail` is called + 3. Device immediately signals to host (Low-Level Runtime) + - Host side: + 1. The assert failure gets detected by Low-Level Runtime + 2. Low-Level Runtime sets event status + 3. Upon call to `sycl::queue::wait_and_throw()` or + `sycl::event::wait_and_throw()` DPCPP Runtime checks event status and + throws "assert" exception + ### Fallback approach @@ -155,6 +176,20 @@ information. DPCPP Runtime checks contents of the assert buffer for assert failure flag after kernel finishes. +The following sequence of events describes how user code gets notified: + - Device side: + 1. Assert fails in device-code in kernel + 2. Fallback version of `__devicelib_assert_fail` is called + 3. Assert information is stored into assert buffer + 4. Kernel continues running + - Host side: + 1. Upon call to `sycl::queue::wait_and_throw()` or + `sycl::event::wait_and_throw()` DPCPP Runtime waits until kernel finishes + and checks assert buffer for assert information throws exception + + +#### Storing accessor metadata and writing assert failure to buffer + Both storing of accessor metadata and writing assert failure is performed with help of built-ins. Implementations of these builtins are substituted by frontend. From a31b808075c2eef9717434863ffc131551bd4c58 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 14 Apr 2021 17:47:38 +0300 Subject: [PATCH 23/45] Add notes on property set usage Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 115 ++++++++++++++++++++++++++++++++++++++++++++- 1 file changed, 114 insertions(+), 1 deletion(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index 8a6767af12c15..e5d0bf2fe76ef 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -188,7 +188,120 @@ The following sequence of events describes how user code gets notified: and checks assert buffer for assert information throws exception -#### Storing accessor metadata and writing assert failure to buffer +#### Online-linking fallback `__devicelib_assert_fail` + +Online linking against fallback implementation of `__devicelib_assert_fail` is +performed only when assertion is enabled. + +In DPCPP headers one can see if assert is enabled with status of `NDEBUG` macro +with `#ifdef`'s. This allows to add implicit buffer argument to kernel +invocation. Here "implicit" means "implicit to the user". + +When in DPCPP Runtime Library this knowledge is obtained from device binary +image descriptor's property sets. + +Each device image is supplied with an array of property sets: +```c++ +struct pi_device_binary_struct { + //... + // Array of property sets + pi_device_binary_property_set PropertySetsBegin; + pi_device_binary_property_set PropertySetsEnd; +}; +``` +Each property set is represented by the following struct: +```c++ +// Named array of properties. +struct _pi_device_binary_property_set_struct { + char *Name; // the name + pi_device_binary_property PropertiesBegin; // array start + pi_device_binary_property PropertiesEnd; // array end +}; +``` +It contains name of property set and array of properties. Each property is +represented by the following struct: +```c++ +struct _pi_device_binary_property_struct { + char *Name; // null-terminated property name + void *ValAddr; // address of property value + uint32_t Type; // _pi_property_type + uint64_t ValSize; // size of property value in bytes +}; +``` + +Whenever `isAssertEnabled` property set is present, this specific device image +was built with `NDEBUG` macro undefined and it requires fallback implementation +of `__devicelib_assert_fail` (i.e. if Device-side Runtime doesn't support it). + +Any properties in `isAssertEnabled` property set are ignored. + +The property set is added to device binary descriptor whenever at least single +translation unit was compiled with assertions enabled i.e. `NDEBUG` undefined. + + +##### Compiling with assert enabled/disabled + +Consider the following two use-case: +```c++ +// impl.cpp +using namespace sycl; +int calculus(int X) { + assert(X && "Invalid value"); + return X * 2; +} + +void enqueueKernel(queue &Q, buffer &B) { + Q.submit([](handler &H) { + auto Acc = B.get_access(H); + H.parallel_for(/* range */, [](item It) { + assert(Acc[It]); + // ... + }); + }); +} + +// main.cpp +// ... +using namespace sycl; + +SYCL_EXTERNAL int calculus(int); +void enqueueKernel(queue&, buffer&); + +void workload() { + queue Q; + buffer B; + + Q.submit([](handler &H) { + auto Acc = B.get_access(H); + H.parallel_for(/* range */, [](item It) { + int X = calculus(0); // should fail assertion + assert(X && "Nil in result"); + Acc[It] = X; + }); + }); + + enqueueKernel(Q, B); + ... +} +``` + +These two files are compiled into a single binary application. There are four +states of definedness of `NDEBUG` macro available: + +| # | `impl.cpp` | `main.cpp` | +| - | ---------- | ---------- | +| 1 | defined | defined | +| 2 | defined | undefined | +| 3 | undefined | defined | +| 4 | undefined | undefined | + +States of definedness of `NDEBUG` macro defines the set of assertions which can +fail. Having assertions enabled in at least one translation unit with device +code requires for `isAssertEnabled` property set being present in device image +descriptor structure. + + +### Storing accessor metadata and writing assert failure to buffer Both storing of accessor metadata and writing assert failure is performed with help of built-ins. Implementations of these builtins are substituted by From 257054a5bfd36c6fad3321a70579602bab16ac8e Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 14 Apr 2021 17:52:31 +0300 Subject: [PATCH 24/45] Address comments Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 15 +++-- .../Assert/SYCL_ONEAPI_ASSERT.asciidoc | 57 ++++++++++++++++--- 2 files changed, 59 insertions(+), 13 deletions(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index e5d0bf2fe76ef..ee5ba258ca9d6 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -14,16 +14,21 @@ compile time. ## Use-case example -``` +```c++ +#include +#include + using namespace sycl; auto ErrorHandler = [] (exception_list Exs) { - for (exception_ptr const& E : Exs) { + for (std::exception_ptr const& E : Exs) { try { std::rethrow_exception(E); } - catch (event_error const& Ex) { - std::cout << “Exception - ” << Ex.what(); // assertion failed - std::abort(); + catch (const exception& Ex) { + if (Ex.code() == errc::ext_oneapi_assert) { + std::cout << “Exception - ” << Ex.what(); // assertion failed + std::abort(); + } } } }; diff --git a/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc b/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc index 8b79bf86fa967..c8a64d2262eed 100644 --- a/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc +++ b/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc @@ -63,11 +63,12 @@ implementation. Catching the asynchronous exception with `sycl::errc::ext_oneapi_assert` error code means that assert failure had happened in device code during kernel -execution and the assert message is already printed to `stderr`. Format of the -assert message is the following: -``` -:: : global id: [,,], local id: [,,] Assertion `` failed. -``` +execution and the assert message is already printed to `stderr`. +The format of the assert message is unspecified, but it will always include the +text of the failing expression, the values of the standard macros `__FILE__` and +`__LINE__`, and the value of the standard variable `__func__`. If the failing +assert comes from an `nd_range` `parallel_for` it will also include the global +ID and the local ID of the failing work item. It is unspecified whether a failing `assert()` returns to its caller before the kernel terminates. If a failing call returns, the device code may need to @@ -82,9 +83,49 @@ assertion. The `assert()` macro is defined in system include headers, not in SYCL headers. On most of systems it is `` and/or `` header files. The user can can disable assertions in device code by defining the `NDEBUG` -preprocessor macro at the point when `` and `/` -are included. - +preprocessor macro prior to including either of `` and +`/`. + +Following is an example use-case: + +#[source] +---- +#include +#include + +using namespace sycl; +auto ErrorHandler = [] (exception_list Exs) { + for (std::exception_ptr const& E : Exs) { + try { + std::rethrow_exception(E); + } + catch (const exception& Ex) { + if (Ex.code() == errc::ext_oneapi_assert) { + std::cout << “Exception - ” << Ex.what(); // assertion failed + std::abort(); + } + } + } +}; + +void user_func(item<2> Item) { + assert( (Item[0] % 2) && “Nil”); +} + +int main() { + queue Q(ErrorHandler); + q.submit([&] (handler& CGH) { + CGH.parallel_for(range<2>{N, M}, [=](item<2> It) { + do_smth(); + user_func(It); + do_smth_else(); + }); + }); + Q.wait_and_throw(); + std::cout << “One shouldn’t see this message.“; + return 0; +} +---- == Feature test macro From 3f501730e1cd89c0760ea3c595200df9d6a34fe1 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 14 Apr 2021 17:55:57 +0300 Subject: [PATCH 25/45] Fix typo and format note Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 11 ++++++----- .../doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc | 2 +- 2 files changed, 7 insertions(+), 6 deletions(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index ee5ba258ca9d6..09042560bf532 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -97,10 +97,11 @@ practical cases. `assert(expr)` macro ends up in call to `__devicelib_assert_fail`. This function is part of [Device library extension](extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst#cl_intel_devicelib_cassert). -Format of assert failure message, printed to `stderr` is the following: -``` -:: : global id: [,,], local id: [,,] Assertion `` failed. -``` +The format of the assert message is unspecified, but it will always include the +text of the failing expression, the values of the standard macros `__FILE__` and +`__LINE__`, and the value of the standard variable `__func__`. If the failing +assert comes from an `nd_range` `parallel_for` it will also include the global +ID and the local ID of the failing work item. Implementation of this function is supplied by Native Device Compiler for safe approach or by DPCPP Compiler for fallback one. @@ -246,7 +247,7 @@ translation unit was compiled with assertions enabled i.e. `NDEBUG` undefined. ##### Compiling with assert enabled/disabled -Consider the following two use-case: +Consider the following example sources: ```c++ // impl.cpp using namespace sycl; diff --git a/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc b/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc index c8a64d2262eed..3ca160ea7d810 100644 --- a/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc +++ b/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc @@ -88,7 +88,7 @@ preprocessor macro prior to including either of `` and Following is an example use-case: -#[source] +[source] ---- #include #include From c1326aa2defe52ca960f3f5efed11d6145db4d2f Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 14 Apr 2021 17:59:34 +0300 Subject: [PATCH 26/45] Fix typo Signed-off-by: Sergey Kanaev --- sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc b/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc index 3ca160ea7d810..284101fe8bac2 100644 --- a/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc +++ b/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc @@ -82,7 +82,7 @@ assertion. The `assert()` macro is defined in system include headers, not in SYCL headers. On most of systems it is `` and/or `` header files. -The user can can disable assertions in device code by defining the `NDEBUG` +The user can disable assertions in device code by defining the `NDEBUG` preprocessor macro prior to including either of `` and `/`. From 5095b1a59fb3e695b9da6fd02849c85ea6cb5f3c Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 14 Apr 2021 18:02:27 +0300 Subject: [PATCH 27/45] Add extension to README Signed-off-by: Sergey Kanaev --- sycl/doc/extensions/README.md | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/doc/extensions/README.md b/sycl/doc/extensions/README.md index e4b4a7bdb52be..239efe8421a9f 100755 --- a/sycl/doc/extensions/README.md +++ b/sycl/doc/extensions/README.md @@ -37,6 +37,7 @@ DPC++ extensions status: | [Use Pinned Memory Property](UsePinnedMemoryProperty/UsePinnedMemoryPropery.adoc) | Supported | | | [Level-Zero backend specification](LevelZeroBackend/LevelZeroBackend.md) | Supported | | | [ITT annotations support](ITTAnnotations/ITTAnnotations.rst) | Supported | | +| [Assert](Assert/SYCL_ONEAPI_ASSERT.asciidoc) | Proposal | | Legend: From 5078fccf940173d6a90829209a503cf3f986e370 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 14 Apr 2021 18:08:08 +0300 Subject: [PATCH 28/45] Note on how property set gets to be set Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 3 +++ 1 file changed, 3 insertions(+) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index 09042560bf532..993f5b5ac681c 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -244,6 +244,9 @@ Any properties in `isAssertEnabled` property set are ignored. The property set is added to device binary descriptor whenever at least single translation unit was compiled with assertions enabled i.e. `NDEBUG` undefined. +The property set is added by `sycl-post-link` tool depending on module metadata. +Metadata is provided by Clang frontend. Metadata name is `is_assert_enabled`. + ##### Compiling with assert enabled/disabled From 9bcac020d074c9cc09d072485770ac73c44fa013 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Thu, 15 Apr 2021 16:18:10 +0300 Subject: [PATCH 29/45] Partially remove mentioning of async exception throw Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 14 +++---- .../Assert/SYCL_ONEAPI_ASSERT.asciidoc | 42 ++++--------------- 2 files changed, 14 insertions(+), 42 deletions(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index 993f5b5ac681c..148db056d6cd4 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -157,10 +157,8 @@ The following sequence of events describes how user code gets notified: 3. Device immediately signals to host (Low-Level Runtime) - Host side: 1. The assert failure gets detected by Low-Level Runtime - 2. Low-Level Runtime sets event status - 3. Upon call to `sycl::queue::wait_and_throw()` or - `sycl::event::wait_and_throw()` DPCPP Runtime checks event status and - throws "assert" exception + 2. Low-Level Runtime prints assert failure message to `stderr` + 3. Low-Level Runtime calls `abort()` ### Fallback approach @@ -189,9 +187,11 @@ The following sequence of events describes how user code gets notified: 3. Assert information is stored into assert buffer 4. Kernel continues running - Host side: - 1. Upon call to `sycl::queue::wait_and_throw()` or - `sycl::event::wait_and_throw()` DPCPP Runtime waits until kernel finishes - and checks assert buffer for assert information throws exception + 1. A distinct thread is launched no later than the point of enqueue of the of + kernel with assertions + 2. This thread polls the enqueued kernels for finish and checks the assert + buffer for assert data + 3. If assert data is present DPCPP Runtime calls `abort()` #### Online-linking fallback `__devicelib_assert_fail` diff --git a/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc b/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc index 284101fe8bac2..9b2b28cce2acb 100644 --- a/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc +++ b/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc @@ -52,18 +52,8 @@ products. This extension adds the ability for device code to call the C++ `assert()` macro. The behavior of `assert()` in device code is similar to its behavior in host code. If the asserted condition is false, a message is printed to `stderr` -and then the program typically aborts. The mechanism for aborting the program is -different, though. Whereas the host version calls `std::abort()`, the device -version causes an asynchronous SYCL `exception` with the -`errc::ext_oneapi_assert` error code to be thrown. The application can catch -this exception like any other asynchronous exception that is thrown from an -executing kernel. The numeric value of this enumeration element is defined by -implementation. - - -Catching the asynchronous exception with `sycl::errc::ext_oneapi_assert` error -code means that assert failure had happened in device code during kernel -execution and the assert message is already printed to `stderr`. +and then the program aborts with `std::abort()`. + The format of the assert message is unspecified, but it will always include the text of the failing expression, the values of the standard macros `__FILE__` and `__LINE__`, and the value of the standard variable `__func__`. If the failing @@ -73,12 +63,7 @@ ID and the local ID of the failing work item. It is unspecified whether a failing `assert()` returns to its caller before the kernel terminates. If a failing call returns, the device code may need to continue execution without deadlocking for the assertion message to be printed -or for the exception to be thrown. - -The contents of the exception's `e.what()` string is unspecified. Since the -assertion message is already printed to `stderr` by the time the exception is -thrown, the `e.what()` string may not include the location of the failing -assertion. +or for `std::abort()` to be called. The `assert()` macro is defined in system include headers, not in SYCL headers. On most of systems it is `` and/or `` header files. @@ -94,34 +79,21 @@ Following is an example use-case: #include using namespace sycl; -auto ErrorHandler = [] (exception_list Exs) { - for (std::exception_ptr const& E : Exs) { - try { - std::rethrow_exception(E); - } - catch (const exception& Ex) { - if (Ex.code() == errc::ext_oneapi_assert) { - std::cout << “Exception - ” << Ex.what(); // assertion failed - std::abort(); - } - } - } -}; void user_func(item<2> Item) { - assert( (Item[0] % 2) && “Nil”); + assert((Item[0] % 2) && “Nil”); } int main() { - queue Q(ErrorHandler); - q.submit([&] (handler& CGH) { + queue Q; + Q.submit([&] (handler& CGH) { CGH.parallel_for(range<2>{N, M}, [=](item<2> It) { do_smth(); user_func(It); do_smth_else(); }); }); - Q.wait_and_throw(); + Q.wait(); std::cout << “One shouldn’t see this message.“; return 0; } From 7ec3ac8676d8f4db10c69f3de6f19ead0591b2ba Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Thu, 15 Apr 2021 17:30:11 +0300 Subject: [PATCH 30/45] Add Assert.md to index Signed-off-by: Sergey Kanaev --- sycl/doc/index.rst | 2 ++ 1 file changed, 2 insertions(+) diff --git a/sycl/doc/index.rst b/sycl/doc/index.rst index 8089d12230730..9be7037fbd959 100644 --- a/sycl/doc/index.rst +++ b/sycl/doc/index.rst @@ -32,3 +32,5 @@ Developing oneAPI DPC++ Compiler KernelProgramCache GlobalObjectsInRuntime LinkedAllocations + Assert + From 8cbfde7a82d845bdb1e16d809060e609c16b4ad9 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Thu, 15 Apr 2021 17:30:31 +0300 Subject: [PATCH 31/45] Remove the rest of exception throws Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 45 +++++++++------------------------------------ 1 file changed, 9 insertions(+), 36 deletions(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index 148db056d6cd4..086cc948d815c 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -5,9 +5,8 @@ Using the standard C++ `assert` API ("assertions") is an important debugging technique widely used by developers. This document describes the design of supporting assertions within SYCL device code. -The basic approach we chose is delivering device-side assertions as host-side -asynchronous exceptions, which allows further extensibility, such as better -error handling or potential recovery. +The basic approach we chose is delivering device-side assertions as call to +`std::abort()` at host-side. As usual, device-side assertions can be disabled by defining `NDEBUG` macro at compile time. @@ -19,42 +18,30 @@ compile time. #include using namespace sycl; -auto ErrorHandler = [] (exception_list Exs) { - for (std::exception_ptr const& E : Exs) { - try { - std::rethrow_exception(E); - } - catch (const exception& Ex) { - if (Ex.code() == errc::ext_oneapi_assert) { - std::cout << “Exception - ” << Ex.what(); // assertion failed - std::abort(); - } - } - } -}; void user_func(item<2> Item) { assert((Item[0] % 2) && “Nil”); } int main() { - queue Q(ErrorHandler); - q.submit([&] (handler& CGH) { + queue Q; + Q.submit([&] (handler& CGH) { CGH.parallel_for(range<2>{N, M}, [=](item<2> It) { do_smth(); user_func(It); do_smth_else(); }); }); - Q.wait_and_throw(); + Q.wait(); std::cout << “One shouldn’t see this message.“; return 0; } ``` In this use-case every work-item with even X dimension will trigger assertion -failure. Assertion failure should be reported via asynchronous exceptions with -[`assert` error code](extensions/Assert/SYCL_INTEL_assert_exception.asciidoc). +failure. Assertion failure should be trigger a call to `std::abort()` at host as +described in +[extension](extensions/Assert/SYCL_INTEL_ASSERT.asciidoc). Even though multiple failures of the same or different assertions can happen in multiple workitems, implementation is required to deliver only one. The assertion failure message is printed to `stderr` by DPCPP Runtime. @@ -69,7 +56,7 @@ From user's point of view there are the following requirements: | # | Title | Description | Importance | | - | ----- | ----------- | ---------- | -| 1 | Handle assertion failure | Signal about assertion failure via SYCL asynchronous exception | Must have | +| 1 | Abort DPC++ application | Abort host application when assert function is called and print a message about assertion | Must have | | 2 | Print assert message | Assert function should print message to stderr at host | Must have | | 3 | Stop under debugger | When debugger is attached, break at assertion point | Highly desired | | 4 | Reliability | Assert failure should be reported regardless of kernel deadlock | Highly desired | @@ -111,20 +98,6 @@ linked against fallback implementation of `__devicelib_assert_fail`. Hence, Native Device Compilers should prefer their implementation instead of the one provided in incoming SPIR-V/LLVM IR binary. -Limitations for user after catching the "assert" asynchronous exception: - - When using GPU device and the kernel hangs/crashes the subsequent enqueues - will fail; - -When using CPU devices the user can proceed with enqueues to the same -device/queue/context. -DPCPP Runtime remains in valid state after "assert" exception been thrown. - - -### Current violation - -While throwing an asynchronous exception is quite an extensible way, for the -time being DPCPP Runtime merely calls `abort()`. - ## Safe approach From cc085f524deb19885c73fb6d355e2235cbb647d4 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Thu, 22 Apr 2021 17:28:22 +0300 Subject: [PATCH 32/45] Address review comments Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 59 ++++++++----------- .../Assert/SYCL_ONEAPI_ASSERT.asciidoc | 8 +-- .../cl_intel_assert_return_code.asciidoc | 4 +- 3 files changed, 32 insertions(+), 39 deletions(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index 086cc948d815c..ebd70246764bb 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -43,8 +43,9 @@ failure. Assertion failure should be trigger a call to `std::abort()` at host as described in [extension](extensions/Assert/SYCL_INTEL_ASSERT.asciidoc). Even though multiple failures of the same or different assertions can happen in -multiple workitems, implementation is required to deliver only one. The -assertion failure message is printed to `stderr` by DPCPP Runtime. +multiple workitems, implementation is required to deliver at least one +assertion. The assertion failure message is printed to `stderr` by DPCPP +Runtime. When multiple kernels are enqueued and more than one fail at assertion, at least single assertion should be reported. @@ -93,10 +94,9 @@ ID and the local ID of the failing work item. Implementation of this function is supplied by Native Device Compiler for safe approach or by DPCPP Compiler for fallback one. -NB: Due to lack of support of online linking in Level-Zero, the application is -linked against fallback implementation of `__devicelib_assert_fail`. Hence, -Native Device Compilers should prefer their implementation instead of the one -provided in incoming SPIR-V/LLVM IR binary. +In order to distinguish which implementation to use, DPCPP Runtime checks for +`cl_intel_devicelib_cassert` extension. If the extension isn't available, then +fallback implementation is used. ## Safe approach @@ -112,15 +112,6 @@ and runtime. The Low-Level Runtime is responsible for: - detecting if assert failure took place; - flushing assert message to `stderr` on host. -When detected, Low-level Runtime reports assert failure to DPCPP Runtime -via events objects. To achieve this, information about assert failure should be -propagated from device-side to SYCL Runtime. This should be performed via calls -to `piEventGetInfo`. This Plugin Interface call "lowers" to `clGetEventInfo` for -OpenCL backend and `zeEventQueryStatus` for Level-Zero backend. - -Refer to [OpenCL](extensions/Assert/opencl.md) and [Level-Zero](extensions/Assert/level-zero.md) -extensions. - The following sequence of events describes how user code gets notified: - Device side: 1. Assert fails in device-code in kernel @@ -142,13 +133,14 @@ Device-side Runtime and Native Device Compiler. Neither it does from Low-level Runtime. Within this approach, a dedicated assert buffer is allocated and implicit kernel -argument is introduced. The argument is an accessor with `discard_read_write` -or `discard_write` access mode. Accessor metadata is stored to program scope -variable. This allows to refer to the accessor without modifying each and every -user's function. Fallback implementation of `__devicelib_assert_fail` restores -accessor metadata from program scope variable and writes assert information to -the assert buffer. Atomic operations are used in order to not overwrite existing -information. +argument is introduced. The argument is an accessor that has either +`access_mode::read_write` or `access_mode::write` access mode and was +constructed with the `property::no_init property`. Accessor metadata is stored +to program scope variable. This allows to refer to the accessor without +modifying each and every user's function. Fallback implementation of +`__devicelib_assert_fail` restores accessor metadata from program scope variable +and writes assert information to the assert buffer. Atomic operations are used +in order to not overwrite existing information. DPCPP Runtime checks contents of the assert buffer for assert failure flag after kernel finishes. @@ -160,8 +152,8 @@ The following sequence of events describes how user code gets notified: 3. Assert information is stored into assert buffer 4. Kernel continues running - Host side: - 1. A distinct thread is launched no later than the point of enqueue of the of - kernel with assertions + 1. A distinct thread is launched no later than the point of enqueue of the + first kernel with assertions 2. This thread polls the enqueued kernels for finish and checks the assert buffer for assert data 3. If assert data is present DPCPP Runtime calls `abort()` @@ -170,7 +162,8 @@ The following sequence of events describes how user code gets notified: #### Online-linking fallback `__devicelib_assert_fail` Online linking against fallback implementation of `__devicelib_assert_fail` is -performed only when assertion is enabled. +performed only when assertion is enabled and Device-side Runtime doesn't provide +implementation of `__devicelib_assert_fail`. In DPCPP headers one can see if assert is enabled with status of `NDEBUG` macro with `#ifdef`'s. This allows to add implicit buffer argument to kernel @@ -208,16 +201,16 @@ struct _pi_device_binary_property_struct { }; ``` -Whenever `isAssertEnabled` property set is present, this specific device image -was built with `NDEBUG` macro undefined and it requires fallback implementation -of `__devicelib_assert_fail` (i.e. if Device-side Runtime doesn't support it). - -Any properties in `isAssertEnabled` property set are ignored. +There's no need for a whole new property set so we reuse `SYCL/misc properties` +property set. Whenever `isAssertEnabled` property is present, this specific +device image was built with `NDEBUG` macro undefined and it requires fallback +implementation of `__devicelib_assert_fail` (i.e. if Device-side Runtime doesn't +support it). -The property set is added to device binary descriptor whenever at least single +The property is added to device binary descriptor whenever at least single translation unit was compiled with assertions enabled i.e. `NDEBUG` undefined. -The property set is added by `sycl-post-link` tool depending on module metadata. +The property is added by `sycl-post-link` tool depending on module metadata. Metadata is provided by Clang frontend. Metadata name is `is_assert_enabled`. @@ -283,7 +276,7 @@ code requires for `isAssertEnabled` property set being present in device image descriptor structure. -### Storing accessor metadata and writing assert failure to buffer +#### Storing accessor metadata and writing assert failure to buffer Both storing of accessor metadata and writing assert failure is performed with help of built-ins. Implementations of these builtins are substituted by diff --git a/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc b/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc index 9b2b28cce2acb..c44d29c72113d 100644 --- a/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc +++ b/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc @@ -55,10 +55,10 @@ host code. If the asserted condition is false, a message is printed to `stderr` and then the program aborts with `std::abort()`. The format of the assert message is unspecified, but it will always include the -text of the failing expression, the values of the standard macros `__FILE__` and -`__LINE__`, and the value of the standard variable `__func__`. If the failing -assert comes from an `nd_range` `parallel_for` it will also include the global -ID and the local ID of the failing work item. +text of the failing expression, the values of the standard macros `+__FILE__+` +and `+__LINE__+`, and the value of the standard variable `+__func__+`. If the +failing assert comes from an `nd_range` `parallel_for` it will also include the +global ID and the local ID of the failing work item. It is unspecified whether a failing `assert()` returns to its caller before the kernel terminates. If a failing call returns, the device code may need to diff --git a/sycl/doc/extensions/Assert/cl_intel_assert_return_code.asciidoc b/sycl/doc/extensions/Assert/cl_intel_assert_return_code.asciidoc index b7eec45d0a26f..58036b0334b4a 100644 --- a/sycl/doc/extensions/Assert/cl_intel_assert_return_code.asciidoc +++ b/sycl/doc/extensions/Assert/cl_intel_assert_return_code.asciidoc @@ -42,9 +42,9 @@ Revision: 1 == Dependencies -This extension is written against the OpenCL Specification Version 1.0, Revision 48. +This extension is written against the OpenCL Specification Version 1.2, Revision 19. -This extension requires OpenCL 1.0 or later. +This extension requires OpenCL 1.2 or later. == Overview From 8835bf8a1ae5dfc1146e455ae672e3e39c1e3caf Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Thu, 6 May 2021 12:05:29 +0300 Subject: [PATCH 33/45] Document program-scope variable approach Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 308 +++++++++++----------- sycl/doc/images/assert-fallback-graph.svg | 3 + 2 files changed, 155 insertions(+), 156 deletions(-) create mode 100644 sycl/doc/images/assert-fallback-graph.svg diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index ebd70246764bb..29073e3e54116 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -76,8 +76,6 @@ practical cases. based on input SPIR-V image. - Low-level Runtime - the backend/runtime behind DPCPP Runtime attached via the Plugin Interface. - - Accessor metadata - parts of accessor representation at device-side: pointer, - ranges, offset. ## How it works? @@ -125,49 +123,73 @@ The following sequence of events describes how user code gets notified: 3. Low-Level Runtime calls `abort()` -### Fallback approach +## Fallback approach -If Device-side Runtime doesn't support `__devicelib_assert_fail` then a buffer -based approach comes in place. The approach doesn't require any support from +If Device-side Runtime doesn't support `__devicelib_assert_fail` then a fallback +approach comes in place. The approach doesn't require any support from Device-side Runtime and Native Device Compiler. Neither it does from Low-level Runtime. -Within this approach, a dedicated assert buffer is allocated and implicit kernel -argument is introduced. The argument is an accessor that has either -`access_mode::read_write` or `access_mode::write` access mode and was -constructed with the `property::no_init property`. Accessor metadata is stored -to program scope variable. This allows to refer to the accessor without -modifying each and every user's function. Fallback implementation of -`__devicelib_assert_fail` restores accessor metadata from program scope variable -and writes assert information to the assert buffer. Atomic operations are used -in order to not overwrite existing information. - -DPCPP Runtime checks contents of the assert buffer for assert failure flag after -kernel finishes. +Within this approach, a mutable program scope variable is introduced. This +variable stores a flag which says if an assert failure was encountered. Fallback +implementation of `__devicelib_assert_fail` atomically raises the flag so that +DPCPP Runtime is able to detect assert failure after kernel finishes. The following sequence of events describes how user code gets notified: - Device side: 1. Assert fails in device-code in kernel 2. Fallback version of `__devicelib_assert_fail` is called - 3. Assert information is stored into assert buffer + 3. Assert information is stored into program-scope variable 4. Kernel continues running - Host side: - 1. A distinct thread is launched no later than the point of enqueue of the - first kernel with assertions - 2. This thread polls the enqueued kernels for finish and checks the assert - buffer for assert data - 3. If assert data is present DPCPP Runtime calls `abort()` + 1. A copy 'kernel' is enqueued as the one depending on user's kernel to get + the value of assert failure flag. + 2. A host-task is enqueued to check value of assert failure flag. + 3. The host task calls abort whenever assert failure flag is set. + +Illustrating this with an example, lets assume the user enqueues three kernels: + - `Kernel #1` + - `Kernel #2` + - `Kernel #3`, which depends on `Kernel #1` + +The resulting graph will look like this: ![graph](images/assert-fallback-graph.svg) + +### Interface to program scope variable + +Multiple translation units could be compiled/linked into a single device binary +image. All of them should have `extern` declaration of program scope variable +available. Definition of the variable is only available within devicelib in the +same binary image where fallback `__devicelib_assert_fail` resides. + +The variable has the following structure and +declaration: + +```c++ +struct AssertHappened { + int Flag = 0; +}; + +#ifdef __SYCL_DEVICE_ONLY__ +extern SYCL_GLOBAL_VAR AssertHappened AssertHappenedMem; +#endif +``` +Here, `SYCL_GLOBAL_VAR` is a macro which wraps special attribute to allow for +mutable program-scope variable. -#### Online-linking fallback `__devicelib_assert_fail` +The reference to extern variable is resolved within online-linking against +fallback devicelib. + +### Online-linking fallback `__devicelib_assert_fail` Online linking against fallback implementation of `__devicelib_assert_fail` is performed only when assertion is enabled and Device-side Runtime doesn't provide implementation of `__devicelib_assert_fail`. In DPCPP headers one can see if assert is enabled with status of `NDEBUG` macro -with `#ifdef`'s. This allows to add implicit buffer argument to kernel -invocation. Here "implicit" means "implicit to the user". +with `#ifdef`'s. This allows to enqueue a copy kernel and host task. The copy +kernel will copy `AssertHappenedMem` to host and host-task will check the `Flag` +value and `abort()` as needed. When in DPCPP Runtime Library this knowledge is obtained from device binary image descriptor's property sets. @@ -213,8 +235,62 @@ translation unit was compiled with assertions enabled i.e. `NDEBUG` undefined. The property is added by `sycl-post-link` tool depending on module metadata. Metadata is provided by Clang frontend. Metadata name is `is_assert_enabled`. +Suppose the following example user code: +```c++ +void user_func(int X) { + assert(X && “X is nil”); +} + +int main() { + queue Q(...); + Q.submit([&] (handler& CGH) { + CGH.single_task([=] () { + do_smth(); + user_func(0); + do_smth_else(); + }); + }); + ... +} +``` + +The following LLVM IR pseudo code will be generated after linking against +fallback implementation of devicelib: +``` +@AssertHappenedMem = global AssertHappened + +/// user's code +void user_func(int X) { +if (!(X && “X is nil")) { + __assert_fail(...); + } +} -##### Compiling with assert enabled/disabled +kernel(...) { + do_smth() + user_func(0); + do_smth_else(); +} + +/// __assert_fail belongs to Linux version of devicelib +void __assert_fail(...) { + ... + __devicelib_assert_fail(...); +} + +void __devicelib_assert_fail(Expr, File, Line, GlobalID, LocalID) { + ... + volatile int *Ptr = (volatile int *)AssertHappenedMem.Flag; + int Expected = 0; + int Desired = 1; + + if (atomic_CAS(&AssertHappenedMem.Flag, Expected, Desired)) + printf("Assertion `%s' failed in %s at line %i. GlobalID: %i, LocalID: %i", + Expr, File, Line, GlobalID, LocalID); +} +``` + +#### Compiling with assert enabled/disabled Consider the following example sources: ```c++ @@ -275,145 +351,65 @@ fail. Having assertions enabled in at least one translation unit with device code requires for `isAssertEnabled` property set being present in device image descriptor structure. +### Raising assert failure flag and reading it on host -#### Storing accessor metadata and writing assert failure to buffer +Each and every translation unit provided by user should have `extern` +declaration of `AssertHappenedMem` i.e. DPCPP headers includes appropriate file +with [declaration](#prog-scope-var-decl). -Both storing of accessor metadata and writing assert failure is performed with -help of built-ins. Implementations of these builtins are substituted by -frontend. +The definition is only provided within devicelib along with +`__devicelib_assert_fail` function which raises the flag. -User's kernel is executed through a wrapper. Wrapping takes place in DPCPP -Runtime headers in a following manner: +Reading of assert failure flag is performed with the help of auxiliary kernel +which is enqueued as dependent on user's one. The flag state is checked later +in host-task. This is achieved with approximately the following changes: -``` -class handler { +```c++ +#include // contains extern decl of AssertHappenedMem -template parallel_for(KernelFunc, Range) { #ifndef NDEBUG - // Assert required -  if (!MQueue->get_device()->assert_fail_supported()) { -    using KName2 = class ASSERT_WRAPPER_NAME(KernelName); -     -    auto AssertBufferAcc = MQueue->get_context()->getAssertBufferAccessor(this); - -    parallel_for_impl( -      Range, -      [=](Item) { -        __store_acc(AssertBuffAcc); -        KernelFunc(Item); -      }); -  } else { +class AssertFlagCopier; #endif - // (No assert required) OR (Assert supported by device) -     // ordinary enqueue process - +class queue { + template event submit(T CGF) { + event Event = submit_impl(CGF); #ifndef NDEBUG - } + // assert required + if (!get_device()->assert_fail_supported()) { + // __devicelib_assert_fail isn't supported by Device-side Runtime + // Linking against fallback impl of __devicelib_assert_fail is performed + // by program manager class + AssertHappened *AH = new AssertHappened; + buffer *Buffer = new buffer{1, AH}; + + // read flag value + event CopierEv = submit_impl([&](handler &CGH) { + CGH.depends_on(Event); + + auto Acc = Buffer->get_access(CGH); + + CGH.single_task([=] { + Acc[0].Flag = atomic_load(&AssertHappenedMem.Flag); + }); + }); + + // check flag state + submit_impl([=](handler &CGH) { + CGH.depends_on(CopierEv); + + CGH.codeplay_host_task([=] { + if (AH->Flag) + abort(); + + free(Buffer); + free(AH); + }); + }); + } #endif -} - -} -``` - - -#### Built-ins operation - -Accessor is a pointer augmented with offset and two ranges (access range and -memory range). - -There are two built-ins provided by frontend: - * `__store_acc()` - to store accessor metadata into program-scope variable. - * `__store_assert_failure()` - to store flag about assert failure in a buffer - using the metadata stored in program-scope variable. - -The accessor should be stored to program scope variable in global address space -using atomic operations. Motivation for using atomic operations: the program may -contain several kernels and some of them could be running simultaneously on a -single device. - -The `__store_assert_failure()` built-in atomically sets a flag in a buffer. The -buffer is accessed using accessor metadata from program-scope variable. This -built-in return a boolean value which is `true` if the flag is set by this call -to `__store_assert_failure()` and `false` if the flag was already set. -Motivation for using atomic operation is the same as with `__store_acc()` -builtin. - -The following pseudo-code snippets shows how these built-ins are used. -First of all, assume the following code as user's one: -``` -void user_func(int X) { - assert(X && “X is nil”); -} - -int main() { - queue Q(...); - Q.submit([&] (handler& CGH) { - CGH.single_task([=] () { - do_smth(); - user_func(0); - do_smth_else(); - }); - }); - ... -} -``` - -The following LLVM IR pseudo code will be generated for the user's code: -``` -@AssertBufferPtr = global void* null -@AssertBufferAccessRange = ... -@AssertBufferMemoryRange = ... -@AssertBufferOffset = ... - -/// user's code -void user_func(int X) { -if (!(X && “X is nil")) { - __assert_fail(...); + return Event; } -} - -users_kernel(...) { - do_smth() - user_func(0); - do_smth_else(); -} - -/// a wrapped user's kernel -kernel(AssertBufferAccessor, OtherArguments...) { - __store_acc(AssertBufferAccessor); - users_kernel(OtherArguments...); -} - -/// __assert_fail belongs to Linux version of devicelib -void __assert_fail(...) { - ... - __devicelib_assert_fail(...); -} - -void __devicelib_assert_fail(Expr, File, Line, GlobalID, LocalID) { - ... - if (__store_assert_info()) - printf("Assertion `%s' failed in %s at line %i. GlobalID: %i, LocalID: %i", - Expr, File, Line, GlobalID, LocalID); -} - -/// The following are built-ins provided by frontend -void __store_acc(accessor) { - %1 = accessor.getPtr(); - store void * %1, void * @AssertBufferPtr -} - -bool __store_assert_info(...) { - AssertBAcc = __fetch_acc(); - // fill in data in AsBAcc - volatile int *Ptr = (volatile int *)AssertBAcc.getPtr(); - bool Expected = false; - bool Desired = true; - - return atomic_cas(Ptr, Expected, Desired, SequentialConsistentMemoryOrder); - // or it could be: - // return !atomic_exchange(Ptr, Desired, SequentialConsistentMemoryOrder); -} +}; ``` diff --git a/sycl/doc/images/assert-fallback-graph.svg b/sycl/doc/images/assert-fallback-graph.svg new file mode 100644 index 0000000000000..fadf4a07ba1c0 --- /dev/null +++ b/sycl/doc/images/assert-fallback-graph.svg @@ -0,0 +1,3 @@ + + +
User's kernel #1
User's kernel #1
User's kernel #2
User's kernel #2
User's kernel #3
User's kernel #3
Copy assert failure flag
Copy assert failure...
Copy assert failure flag
Copy assert failure...
Host-task with check for the value of assert failure flag
Host-task with check...
Host-task with check for the value of assert failure flag
Host-task with check...
Copy assert failure flag
Copy assert failure...
Host-task with check for the value of assert failure flag
Host-task with check...Viewer does not support full SVG 1.1 From ecb8659af546350d6da8f1b19d451c1bb95c3008 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Fri, 7 May 2021 17:23:04 +0300 Subject: [PATCH 34/45] Remove L0 and OCL extensions. Signed-off-by: Sergey Kanaev --- .../cl_intel_assert_return_code.asciidoc | 99 ---------------- .../ze_intel_assert_return_code.asciidoc | 111 ------------------ 2 files changed, 210 deletions(-) delete mode 100644 sycl/doc/extensions/Assert/cl_intel_assert_return_code.asciidoc delete mode 100644 sycl/doc/extensions/Assert/ze_intel_assert_return_code.asciidoc diff --git a/sycl/doc/extensions/Assert/cl_intel_assert_return_code.asciidoc b/sycl/doc/extensions/Assert/cl_intel_assert_return_code.asciidoc deleted file mode 100644 index 58036b0334b4a..0000000000000 --- a/sycl/doc/extensions/Assert/cl_intel_assert_return_code.asciidoc +++ /dev/null @@ -1,99 +0,0 @@ -cl_intel_assert_return_code -====================================== - -// This section needs to be after the document title. -:doctype: book -:toc2: -:toc: left -:encoding: utf-8 -:lang: en - -:blank: pass:[ +] - -// Set the default source code type in this document to C++, -// for syntax highlighting purposes. This is needed because -// docbook uses c++ and html5 uses cpp. -:language: {basebackend@docbook:c++:cpp} - -== Name Strings - -+cl_intel_assert_return_code+ - -== Notice - -Copyright (c) 2021 Intel Corporation. All rights reserved. - -== Status - -Working Draft - -This is a preview extension specification, intended to provide early access to -a feature for review and community feedback. When the feature matures, this -specification may be released as a formal extension. - -Because the interfaces defined by this specification are not final and are -subject to change they are not intended to be used by shipping software -products. - -== Version - -Built On: {docdate} + -Revision: 1 - -== Dependencies - -This extension is written against the OpenCL Specification Version 1.2, Revision 19. - -This extension requires OpenCL 1.2 or later. - -== Overview - -This extension allows OpenCL 1.x and 2.x devices to notify host that assert had -happened. - -== New error code - -[source] ----- -CL_ASSERT_FAILURE ----- - -Negative value of this error code should be set into `param_value` of -`clGetEventInfo` as described in table 5.15 "clGetEventInfo prameter queries" if -assert failure took place in device-code during kernel execution. - -An example: -[source] ----- -cl_event Event; // describes an event of kernel been submitted previously -cl_int Result; -size_t ResultSize; - -clGetEventInfo(Event, CL_EVENT_COMMAND_EXECUTION_STATUS, sizeof(Result), &Result, &ResultSize); ----- - -If kernel failed an assertion `clGetEventInfo` should put `CL_ASSERT_FAILURE` in -`Result`. - -== Issues - -None. - -== Revision History - -[cols="5,15,15,70"] -[grid="rows"] -[options="header"] -|======================================== -|Rev|Date|Author|Changes -|1|2021-04-09|Sergey Kanaev|*Initial public working draft* -|======================================== - -//************************************************************************ -//Other formatting suggestions: -// -//* Use *bold* text for host APIs, or [source] syntax highlighting. -//* Use +mono+ text for device APIs, or [source] syntax highlighting. -//* Use +mono+ text for extension names, types, or enum values. -//* Use _italics_ for parameters. -//************************************************************************ diff --git a/sycl/doc/extensions/Assert/ze_intel_assert_return_code.asciidoc b/sycl/doc/extensions/Assert/ze_intel_assert_return_code.asciidoc deleted file mode 100644 index 85db891ad5bcf..0000000000000 --- a/sycl/doc/extensions/Assert/ze_intel_assert_return_code.asciidoc +++ /dev/null @@ -1,111 +0,0 @@ -ze_intel_assert_return_code -====================================== - -// This section needs to be after the document title. -:doctype: book -:toc2: -:toc: left -:encoding: utf-8 -:lang: en - -:blank: pass:[ +] - -// Set the default source code type in this document to C++, -// for syntax highlighting purposes. This is needed because -// docbook uses c++ and html5 uses cpp. -:language: {basebackend@docbook:c++:cpp} - -== Name Strings - -+ze_intel_assert_return_code+ - -== Notice - -Copyright (c) 2021 Intel Corporation. All rights reserved. - -== Status - -Working Draft - -This is a preview extension specification, intended to provide early access to -a feature for review and community feedback. When the feature matures, this -specification may be released as a formal extension. - -Because the interfaces defined by this specification are not final and are -subject to change they are not intended to be used by shipping software -products. - -== Version - -Built On: {docdate} + -Revision: 1 - -== Dependencies - -This extension is written against the Level-Zero Specification Version 1.1.2. - -== Overview - -This extension allows Level-Zero devices to notify host that assert had -happened. - -== New enumeration value - -`ze_result_t`: + -[source] ----- -ZE_RESULT_ASSERT_FAILED ----- - -This value should be returned by `zeEventQueryStatus` if assert failure took -place in device-code during kernel execution. - -An example: -[source] ----- -ze_event_handle_t Event; // describes an event of kernel been submitted previously -ze_result Result = zeEventQueryStatus(Event); ----- - -If kernel failed an assertion `zeEventQueryStatus` should return -`ZE_RESULT_ASSERT_FAILED`. - - -== Modifications to Level-Zero API - -(Add to Section API Documentation / Core API / Common / Common Enums / `ze_result_t`) :: -+ --- -`ZE_RESULT_ASSERT_FAILED = 0x70000006` + -[Core] Assert failure took place in device-code during kernel execution. --- - -(Add to section API Documentation / Core API / Event / Event Functions / `zeEventQueryStatus`) :: -+ --- -Return: + -`ZE_RESULT_ASSERT_FAILED` --- - -== Issues - -None. - -== Revision History - -[cols="5,15,15,70"] -[grid="rows"] -[options="header"] -|======================================== -|Rev|Date|Author|Changes -|1|2021-04-09|Sergey Kanaev|*Initial public working draft* -|======================================== - -//************************************************************************ -//Other formatting suggestions: -// -//* Use *bold* text for host APIs, or [source] syntax highlighting. -//* Use +mono+ text for device APIs, or [source] syntax highlighting. -//* Use +mono+ text for extension names, types, or enum values. -//* Use _italics_ for parameters. -//************************************************************************ From 07debdb03835d1c3f6c70656c18c7e74a91478a7 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Tue, 11 May 2021 21:42:00 +0300 Subject: [PATCH 35/45] Address comments Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 34 +++++++++++++++++++++------------- 1 file changed, 21 insertions(+), 13 deletions(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index 29073e3e54116..31b86c8d7caa4 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -223,17 +223,20 @@ struct _pi_device_binary_property_struct { }; ``` -There's no need for a whole new property set so we reuse `SYCL/misc properties` -property set. Whenever `isAssertEnabled` property is present, this specific -device image was built with `NDEBUG` macro undefined and it requires fallback -implementation of `__devicelib_assert_fail` (i.e. if Device-side Runtime doesn't -support it). +A distinct property set `SYCL/assert used` is added. In this set a single +with the name of the kernel is added whenever the kernel uses assert. Use of +assert is detected through call to `__devicelib_assert_fail` function after +linking device binary image with wrapper device library (the `libsycl-crt` +library). -The property is added to device binary descriptor whenever at least single -translation unit was compiled with assertions enabled i.e. `NDEBUG` undefined. +The property set and the underlying properties are added by `sycl-post-link` +tool with help of building callgraph for each and every kernel in device binary +image. -The property is added by `sycl-post-link` tool depending on module metadata. -Metadata is provided by Clang frontend. Metadata name is `is_assert_enabled`. +The added property is used for: + - deciding if online-linking against fallback devicelib is required; + - if there's a need to enqueue program scope variable copier kernel and checker + host-task. Suppose the following example user code: ```c++ @@ -353,9 +356,13 @@ descriptor structure. ### Raising assert failure flag and reading it on host -Each and every translation unit provided by user should have `extern` -declaration of `AssertHappenedMem` i.e. DPCPP headers includes appropriate file -with [declaration](#prog-scope-var-decl). +Each and every translation unit provided by user should have declaration of +assert flag read function: +```c++ +int __devicelib_assert_read(void); +``` +Also, the [AssertHappened](#prog-scope-var-decl) structure type should be +available for the copier kernel. The definition is only provided within devicelib along with `__devicelib_assert_fail` function which raises the flag. @@ -375,8 +382,9 @@ class queue { template event submit(T CGF) { event Event = submit_impl(CGF); #ifndef NDEBUG + std::string KernelName = /* get kernel name from calls to parallel_for, etc. */; // assert required - if (!get_device()->assert_fail_supported()) { + if (!get_device()->assert_fail_supported() && isAssertUsed(KernelName)) { // __devicelib_assert_fail isn't supported by Device-side Runtime // Linking against fallback impl of __devicelib_assert_fail is performed // by program manager class From 995e4d8b58f801fa79dc2757d332be4541fe6f9a Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 12 May 2021 10:15:16 +0300 Subject: [PATCH 36/45] Fix typo Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index 31b86c8d7caa4..55f68f9abb6ca 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -223,7 +223,7 @@ struct _pi_device_binary_property_struct { }; ``` -A distinct property set `SYCL/assert used` is added. In this set a single +A distinct property set `SYCL/assert used` is added. In this set a property with the name of the kernel is added whenever the kernel uses assert. Use of assert is detected through call to `__devicelib_assert_fail` function after linking device binary image with wrapper device library (the `libsycl-crt` From b57ac48f17d20127b6a87ce53f12008f6013a2c3 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 12 May 2021 10:41:50 +0300 Subject: [PATCH 37/45] Fix typo Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index 55f68f9abb6ca..4e205338e2f57 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -350,14 +350,12 @@ states of definedness of `NDEBUG` macro available: | 4 | undefined | undefined | States of definedness of `NDEBUG` macro defines the set of assertions which can -fail. Having assertions enabled in at least one translation unit with device -code requires for `isAssertEnabled` property set being present in device image -descriptor structure. +fail. ### Raising assert failure flag and reading it on host Each and every translation unit provided by user should have declaration of -assert flag read function: +assert flag read function available: ```c++ int __devicelib_assert_read(void); ``` From d2f13ff0495991efedefe843ea801a94f78cbcec Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Mon, 17 May 2021 13:37:42 +0300 Subject: [PATCH 38/45] Address review comments Signed-off-by: Sergey Kanaev Co-authored-by: kbobrovs --- sycl/doc/Assert.md | 95 +++++++++++++++++++++------------------------- 1 file changed, 44 insertions(+), 51 deletions(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index 4e205338e2f57..dc9126cc073d7 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -39,13 +39,13 @@ int main() { ``` In this use-case every work-item with even X dimension will trigger assertion -failure. Assertion failure should be trigger a call to `std::abort()` at host as +failure. Assertion failure should trigger a call to `std::abort()` at host as described in [extension](extensions/Assert/SYCL_INTEL_ASSERT.asciidoc). Even though multiple failures of the same or different assertions can happen in multiple workitems, implementation is required to deliver at least one assertion. The assertion failure message is printed to `stderr` by DPCPP -Runtime. +Runtime or underlying backend. When multiple kernels are enqueued and more than one fail at assertion, at least single assertion should be reported. @@ -93,7 +93,7 @@ Implementation of this function is supplied by Native Device Compiler for safe approach or by DPCPP Compiler for fallback one. In order to distinguish which implementation to use, DPCPP Runtime checks for -`cl_intel_devicelib_cassert` extension. If the extension isn't available, then +`PI_INTEL_DEVICELIB_CASSERT` extension. If the extension isn't available, then fallback implementation is used. @@ -101,7 +101,9 @@ fallback implementation is used. This is the preferred approach and implementations should use it when possible. It guarantees assertion failure notification delivery to the host regardless of -kernel behavior which hit the assertion. +kernel behavior which hit the assertion. If backend suports the safe approach, +it must report this capability to DPCPP Runtime via the +`PI_INTEL_DEVICELIB_CASSERT` extension query. The Native Device Compiler is responsible for providing implementation of `__devicelib_assert_fail` which completely hides details of communication @@ -125,10 +127,10 @@ The following sequence of events describes how user code gets notified: ## Fallback approach -If Device-side Runtime doesn't support `__devicelib_assert_fail` then a fallback -approach comes in place. The approach doesn't require any support from -Device-side Runtime and Native Device Compiler. Neither it does from Low-level -Runtime. +If Device-side Runtime doesn't support `__devicelib_assert_fail` (as reported +via `PI_INTEL_DEVICELIB_CASSERT` extension query) then a fallback approach comes +in place. The approach doesn't require any support from Device-side Runtime and +Native Device Compiler. Neither it does from Low-level Runtime. Within this approach, a mutable program scope variable is introduced. This variable stores a flag which says if an assert failure was encountered. Fallback @@ -147,10 +149,15 @@ The following sequence of events describes how user code gets notified: 2. A host-task is enqueued to check value of assert failure flag. 3. The host task calls abort whenever assert failure flag is set. +DPCPP Runtime will automatically check if assertions are enabled in the kernel +being run, and won't enqueue the auxiliary kernels if assertions are not +enabled. So there is no host-side runtime overhead when assertion are not +enabled. + Illustrating this with an example, lets assume the user enqueues three kernels: - - `Kernel #1` - - `Kernel #2` - - `Kernel #3`, which depends on `Kernel #1` + - `Kernel #1`, uses assert + - `Kernel #2`, uses assert + - `Kernel #3`, uses assert and depends on `Kernel #1` The resulting graph will look like this: ![graph](images/assert-fallback-graph.svg) @@ -165,9 +172,15 @@ same binary image where fallback `__devicelib_assert_fail` resides. declaration: ```c++ +namespace cl { +namespace sycl { +namespace detail { struct AssertHappened { int Flag = 0; }; +} +} +} #ifdef __SYCL_DEVICE_ONLY__ extern SYCL_GLOBAL_VAR AssertHappened AssertHappenedMem; @@ -189,49 +202,29 @@ implementation of `__devicelib_assert_fail`. In DPCPP headers one can see if assert is enabled with status of `NDEBUG` macro with `#ifdef`'s. This allows to enqueue a copy kernel and host task. The copy kernel will copy `AssertHappenedMem` to host and host-task will check the `Flag` -value and `abort()` as needed. +value and `abort()` as needed. The kernel and host task are enqueued when +`NDEBUG` macro isn't defined. When in DPCPP Runtime Library this knowledge is obtained from device binary image descriptor's property sets. -Each device image is supplied with an array of property sets: -```c++ -struct pi_device_binary_struct { - //... - // Array of property sets - pi_device_binary_property_set PropertySetsBegin; - pi_device_binary_property_set PropertySetsEnd; -}; -``` -Each property set is represented by the following struct: -```c++ -// Named array of properties. -struct _pi_device_binary_property_set_struct { - char *Name; // the name - pi_device_binary_property PropertiesBegin; // array start - pi_device_binary_property PropertiesEnd; // array end -}; -``` -It contains name of property set and array of properties. Each property is -represented by the following struct: -```c++ -struct _pi_device_binary_property_struct { - char *Name; // null-terminated property name - void *ValAddr; // address of property value - uint32_t Type; // _pi_property_type - uint64_t ValSize; // size of property value in bytes -}; -``` +Each device image is supplied with an array of property sets. For description +of property sets see `struct pi_device_binary_struct` in +[`pi.h`](https://github.com/intel/llvm/blob/sycl/sycl/include/CL/sycl/detail/pi.h#L692) A distinct property set `SYCL/assert used` is added. In this set a property -with the name of the kernel is added whenever the kernel uses assert. Use of -assert is detected through call to `__devicelib_assert_fail` function after -linking device binary image with wrapper device library (the `libsycl-crt` -library). - -The property set and the underlying properties are added by `sycl-post-link` -tool with help of building callgraph for each and every kernel in device binary -image. +with the name of the kernel is added whenever the kernel uses assert. The use of +assert is detected by a specific LLVM IR pass invoked by the `sycl-post-link` +tool which runs on linked device code, i.e. after linking with the `libsycl-crt` +library which defines the assert function. The pass builds complete call graph +for a kernel, and sees if there's a call to `__devicelib_assert_fail` anywhere +in the graph. If found, `sycl-post-link` adds the property for the kernel. + +The same is done for all indirect callable functions (marked with specific +attribute) found in the linked device code. Those are functions whose pointers +can be taken and passed around in device code. If a callgraph for any such +function has a call to `__devicelib_assert_fail`, then all kernels in the module +are conservatively marked as using asserts. The added property is used for: - deciding if online-linking against fallback devicelib is required; @@ -340,7 +333,7 @@ void workload() { ``` These two files are compiled into a single binary application. There are four -states of definedness of `NDEBUG` macro available: +states of definition of `NDEBUG` macro available: | # | `impl.cpp` | `main.cpp` | | - | ---------- | ---------- | @@ -349,12 +342,12 @@ states of definedness of `NDEBUG` macro available: | 3 | undefined | defined | | 4 | undefined | undefined | -States of definedness of `NDEBUG` macro defines the set of assertions which can +States of definition of `NDEBUG` macro defines the set of assertions which can fail. ### Raising assert failure flag and reading it on host -Each and every translation unit provided by user should have declaration of +All translation units provided by the user should have a declaration of the assert flag read function available: ```c++ int __devicelib_assert_read(void); From 6281bc52eb11582ec6a9768f80f5a053cf676368 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 19 May 2021 16:35:10 +0300 Subject: [PATCH 39/45] Switch to __devicelib_assert_read Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index dc9126cc073d7..69ceaa0b40c3e 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -363,10 +363,11 @@ which is enqueued as dependent on user's one. The flag state is checked later in host-task. This is achieved with approximately the following changes: ```c++ -#include // contains extern decl of AssertHappenedMem - #ifndef NDEBUG class AssertFlagCopier; +#ifdef __SYCL_DEVICE_ONLY__ +int __devicelib_assert_read(void); +#endif #endif class queue { @@ -389,7 +390,9 @@ class queue { auto Acc = Buffer->get_access(CGH); CGH.single_task([=] { - Acc[0].Flag = atomic_load(&AssertHappenedMem.Flag); +#ifdef __SYCL_DEVICE_ONLY__ + Acc[0].Flag = __devicelib_assert_read(); +#endif }); }); From a5461f3d5dd86e1d740c40d6cc7cd8d4d922b385 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 19 May 2021 16:35:43 +0300 Subject: [PATCH 40/45] Remove use of NDEBUG from suggested changes Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 4 ---- 1 file changed, 4 deletions(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index 69ceaa0b40c3e..56c384749b3a0 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -363,17 +363,14 @@ which is enqueued as dependent on user's one. The flag state is checked later in host-task. This is achieved with approximately the following changes: ```c++ -#ifndef NDEBUG class AssertFlagCopier; #ifdef __SYCL_DEVICE_ONLY__ int __devicelib_assert_read(void); #endif -#endif class queue { template event submit(T CGF) { event Event = submit_impl(CGF); -#ifndef NDEBUG std::string KernelName = /* get kernel name from calls to parallel_for, etc. */; // assert required if (!get_device()->assert_fail_supported() && isAssertUsed(KernelName)) { @@ -409,7 +406,6 @@ class queue { }); }); } -#endif return Event; } }; From 32a32f46e58fa106c683ec21056a46dfc05d4c1d Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Wed, 19 May 2021 17:32:33 +0300 Subject: [PATCH 41/45] Reorder text to increase readability Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 22 +++++++++++++++------- 1 file changed, 15 insertions(+), 7 deletions(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index 56c384749b3a0..1a397674f49cf 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -200,13 +200,8 @@ performed only when assertion is enabled and Device-side Runtime doesn't provide implementation of `__devicelib_assert_fail`. In DPCPP headers one can see if assert is enabled with status of `NDEBUG` macro -with `#ifdef`'s. This allows to enqueue a copy kernel and host task. The copy -kernel will copy `AssertHappenedMem` to host and host-task will check the `Flag` -value and `abort()` as needed. The kernel and host task are enqueued when -`NDEBUG` macro isn't defined. - -When in DPCPP Runtime Library this knowledge is obtained from device binary -image descriptor's property sets. +with `#ifdef`'s. When in DPCPP Runtime Library this knowledge is obtained from +device binary image descriptor's property sets. Each device image is supplied with an array of property sets. For description of property sets see `struct pi_device_binary_struct` in @@ -347,6 +342,19 @@ fail. ### Raising assert failure flag and reading it on host +In DPCPP headers one can see if assert is enabled with status of `NDEBUG` macro +with `#ifdef`'s. Though, in order to support for multi translation unit use-case +it's not allowed to rely on definition of `NDEBUG` macro. + +*Note: Multi translation unit use-case here is the one with `SYCL_EXTERNAL` +function compiled with assertions enabled and used in a kernel but the kernel +is compiled with assertions disabled.* + +There're two commands used for reading assert failure flag: copy kernel and +checker host task. The copy kernel will copy `AssertHappenedMem` to host and +host-task will check the `Flag` value and `abort()` as needed. The kernel and +host task are enqueued when `NDEBUG` macro isn't defined. + All translation units provided by the user should have a declaration of the assert flag read function available: ```c++ From 641d07178f240f1b3f4d63e9f1836095f274eeae Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Thu, 20 May 2021 11:58:27 +0300 Subject: [PATCH 42/45] Address review comment Signed-off-by: Sergey Kanaev Co-authored-by: kbobrovs --- sycl/doc/Assert.md | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index 1a397674f49cf..972352450c45f 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -353,7 +353,9 @@ is compiled with assertions disabled.* There're two commands used for reading assert failure flag: copy kernel and checker host task. The copy kernel will copy `AssertHappenedMem` to host and host-task will check the `Flag` value and `abort()` as needed. The kernel and -host task are enqueued when `NDEBUG` macro isn't defined. +host task are enqueued together with a kernel only when the corresponding device +binary image for this kernel tells that it may use (maybe indirectly) the +`assert` in its code. All translation units provided by the user should have a declaration of the assert flag read function available: From dc058a9fe3c1deb7e51fe5cc1363f0cc489c1f9d Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Thu, 27 May 2021 13:23:09 +0300 Subject: [PATCH 43/45] Address review comments Co-authored-by: bader Signed-off-by: Sergey Kanaev --- sycl/doc/Assert.md | 10 +++++----- sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc | 4 +++- 2 files changed, 8 insertions(+), 6 deletions(-) diff --git a/sycl/doc/Assert.md b/sycl/doc/Assert.md index 972352450c45f..12b074c258665 100644 --- a/sycl/doc/Assert.md +++ b/sycl/doc/Assert.md @@ -38,17 +38,17 @@ int main() { } ``` -In this use-case every work-item with even X dimension will trigger assertion -failure. Assertion failure should trigger a call to `std::abort()` at host as -described in +In this use-case every work-item with even index along 0 dimension will trigger +assertion failure. Assertion failure should trigger a call to `std::abort()` at +host as described in [extension](extensions/Assert/SYCL_INTEL_ASSERT.asciidoc). Even though multiple failures of the same or different assertions can happen in -multiple workitems, implementation is required to deliver at least one +multiple work-items, implementation is required to deliver at least one assertion. The assertion failure message is printed to `stderr` by DPCPP Runtime or underlying backend. When multiple kernels are enqueued and more than one fail at assertion, at least -single assertion should be reported. +one assertion should be reported. ## User requirements diff --git a/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc b/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc index c44d29c72113d..b2ab21ea04561 100644 --- a/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc +++ b/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc @@ -21,9 +21,10 @@ :cpp: C++ == Notice + IMPORTANT: This specification is a draft. -Copyright (c) 2021-2021 Intel Corporation. All rights reserved. +Copyright (c) 2021 Intel Corporation. All rights reserved. NOTE: Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks of The Khronos Group Inc. OpenCL(TM) is a trademark of Apple Inc. @@ -49,6 +50,7 @@ subject to change they are not intended to be used by shipping software products. == Introduction + This extension adds the ability for device code to call the C++ `assert()` macro. The behavior of `assert()` in device code is similar to its behavior in host code. If the asserted condition is false, a message is printed to `stderr` From 16fd8f0e5ab18626cb5860ce1229e047a6816cf1 Mon Sep 17 00:00:00 2001 From: Sergey Kanaev Date: Thu, 27 May 2021 17:56:33 +0300 Subject: [PATCH 44/45] Add aspect Signed-off-by: Sergey Kanaev --- .../Assert/SYCL_ONEAPI_ASSERT.asciidoc | 16 ++++++++++++++++ 1 file changed, 16 insertions(+) diff --git a/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc b/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc index b2ab21ea04561..dffe35bab9958 100644 --- a/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc +++ b/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc @@ -117,6 +117,22 @@ extension’s APIs the implementation supports. |1 |Initial extension version. Base features are supported. |=== +== Extension to `enum class aspect` + +[source] +---- +namespace sycl { +enum class aspect { + ext_oneapi_native_assert +} +} +---- + +If device has the `ext_oneapi_native_assert` aspect, then its Device-Side +Runtime is capable of native support of `assert`. That is, safe implementation +is used. If device doesn't have the aspect, then fallback implementation is +used. + == Version Built On: {docdate} + From fbca768d88526d9a949f147d2c42b7e067908fd1 Mon Sep 17 00:00:00 2001 From: sergei <57672082+s-kanaev@users.noreply.github.com> Date: Thu, 27 May 2021 20:48:29 +0300 Subject: [PATCH 45/45] Update extension with suggestion Signed-off-by: Sergey Kanaev --- .../Assert/SYCL_ONEAPI_ASSERT.asciidoc | 16 ++++++++++++---- 1 file changed, 12 insertions(+), 4 deletions(-) diff --git a/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc b/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc index dffe35bab9958..24004b525d37d 100644 --- a/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc +++ b/sycl/doc/extensions/Assert/SYCL_ONEAPI_ASSERT.asciidoc @@ -62,10 +62,18 @@ and `+__LINE__+`, and the value of the standard variable `+__func__+`. If the failing assert comes from an `nd_range` `parallel_for` it will also include the global ID and the local ID of the failing work item. -It is unspecified whether a failing `assert()` returns to its caller before the -kernel terminates. If a failing call returns, the device code may need to -continue execution without deadlocking for the assertion message to be printed -or for `std::abort()` to be called. +Some devices implement `assert()` natively while others use a fallback +implementation, and the two implementations provide different guarantees. The +native implementation is most similar to the way `assert()` works on the host. If +an assertion fails in the native implementation, the assertion message is +immediately printed to stderr and the program terminates by calling +`std::abort()`. If an assertion fails with the fallback implementation, the +failing assert() returns back to its caller and the device code must continue +executing (without deadlocking) until the kernel completes. The implementation +prints the assertion message to stderr and terminates with `std::abort()` only +after the kernel completes execution. An application can determine which of the +two mechanisms a device uses by testing the device aspect +`aspect::ext_oneapi_native_assert`. The `assert()` macro is defined in system include headers, not in SYCL headers. On most of systems it is `` and/or `` header files.