|
Important
|
This specification is a draft. |
Copyright (c) 2021 Intel Corporation. All rights reserved.
|
Note
|
Khronos® is a registered trademark and SYCL™ and SPIR™ are trademarks of The Khronos Group Inc. OpenCL™ 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. |
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.
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 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.
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 <cassert> and/or <assert.h> header files.
The user can disable assertions in device code by defining the NDEBUG
preprocessor macro prior to including either of <sycl.hpp> and
<cassert>/<assert.h>.
Following is an example use-case:
#include <cassert>
#include <sycl/sycl.hpp>
using namespace sycl;
void user_func(item<2> Item) {
assert((Item[0] % 2) && “Nil”);
}
int main() {
queue Q;
Q.submit([&] (handler& CGH) {
CGH.parallel_for<class TheKernel>(range<2>{N, M}, [=](item<2> It) {
do_smth();
user_func(It);
do_smth_else();
});
});
Q.wait();
std::cout << “One shouldn’t see this message.“;
return 0;
}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.
| Value | Description |
|---|---|
1 |
Initial extension version. Base features are supported. |
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.