SYCL Support

Description

All integer types of this library (unsigned u8-u128, signed i8-i128, and bounded types) and the safe floating-point types f32 and f64 support compilation with the Intel oneAPI DPC++ compiler (icpx) and use inside SYCL device kernels.

SYCL support is opt-in. Define BOOST_SAFE_NUMBERS_ENABLE_SYCL and include <sycl/sycl.hpp> before any Boost.safe_numbers header, so that the SYCL_EXTERNAL keyword is available:

#include <sycl/sycl.hpp>                              // must come first
#include <boost/safe_numbers/unsigned_integers.hpp>
#include <boost/safe_numbers/device_error_reporting.hpp>

Build with icpx -fsycl -DBOOST_SAFE_NUMBERS_ENABLE_SYCL=1.

Unsigned multiplication of the small integer types (u8, u16, u32) lowers to the llvm.umul.with.overflow intrinsic. The Intel oneAPI DPC++ range-rounding wrapper that a plain parallel_for over a non-round global range generates cannot materialize that intrinsic and aborts the kernel at JIT time. If you multiply small unsigned safe types inside such a kernel, either launch over an nd_range (which does not use the wrapper) or compile with -DSYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING=1. The library’s own SYCL tests use the latter.

Because a SYCL device kernel cannot throw, an error that would raise an exception on the host is instead recorded on the device and re-raised on the host by a device_error_context. Unlike a raw kernel launch, where you would call queue::wait() yourself, you wrap the launch in a context:

using test_type = boost::safe_numbers::u128;

int main()
{
    sycl::queue q;

    // Setup: allocate shared memory, generate inputs, etc.

    // The context binds to the queue used for the kernels. On synchronize() it
    // throws the same exception type the operation would have thrown on the host.
    boost::safe_numbers::device_error_context ctx {q};

    q.parallel_for(sycl::range<1>(num_elements), [=](sycl::id<1> idx)
    {
        const int i {static_cast<int>(idx[0])};
        out[i] = boost::safe_numbers::bit_ceil(in[i]);
    });

    try
    {
        // Waits for all enqueued kernels, then throws if the device recorded an error.
        ctx.synchronize();
    }
    catch (const /*exception type*/& e)
    {
        // Perform error handling
    }
}

An exception thrown by ctx.synchronize() will have an e.what() that looks something like:

Device error on thread 1234 at /path/to/boost/safe_numbers/include/boost/safe_numbers/detail/unsigned_integer_basis.hpp:1067: Underflow detected in u16 subtraction

The reported thread id is the failing work-item’s dimension-0 global linear id (the analog of the CUDA backend’s blockIdx.x * blockDim.x + threadIdx.x), so a single bad input element can be narrowed down exactly as on CUDA. It is populated by default whenever the implementation provides the free-function queries extension (feature-test macro SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES, present in DPC++), for plain range kernels, nd_range kernels, and single_task (which reports 0) alike; kernels launched over more than one dimension report only dimension 0. Without the extension, or with BOOST_SAFE_NUMBERS_SYCL_DISABLE_WORKITEM_ID defined, the id is -1 ("unknown").

The device-side error buffer is a device_global, provided by the sycl_ext_oneapi_device_global extension (feature-test macro SYCL_EXT_ONEAPI_DEVICE_GLOBAL). Core SYCL 2020 forbids mutable program-scope variables in device code, so an extension is the only way a deeply nested operation can reach the shared error record implicitly, i.e. without changing how you write arithmetic in kernels. The extension is available in DPC++ (icpx) across its device backends; on an implementation that does not advertise it, including <boost/safe_numbers/device_error_reporting.hpp> under BOOST_SAFE_NUMBERS_ENABLE_SYCL produces a clear compile-time error. The work-item id in the error record similarly comes from the sycl_ext_oneapi_free_function_queries extension, which is optional: without it the id degrades to -1 instead of failing the build. The safe types themselves (and the value-based checked_*/overflowing_*/saturating_* functions, which report errors through their return values instead) remain usable on any conforming SYCL 2020 implementation.

Floating-Point Support

The safe floating-point types f32 and f64 compile with icpx and execute inside device kernels alongside the integer types. Their four checked operators (+, -, *, and /) run on the device and classify the result according to IEEE 754 exactly as they do on the host. As on the host there is no operator%; see Floating-Point Types.

When a device-side floating-point operation produces an exceptional result, it records the error into the device_error_context just as the integer operations do, and synchronize() throws the matching exception on the host:

Device Floating-Point Result Host Exception

Saturation to positive infinity

std::overflow_error

Saturation to negative infinity

std::underflow_error

Invalid operation, a NaN operand, or division of a finite non-zero value by zero

std::domain_error

The range-checked bounded_float type runs on the device as well, as do the bounded integer types bounded_uint and bounded_int; a result that leaves the declared [Min, Max] range is captured by the same device_error_context and rethrown on the host.

The device_exception_mode Enum

#include <boost/safe_numbers/device_error_reporting.hpp>

namespace boost::safe_numbers {

enum class device_exception_mode : unsigned
{
    trapped,
    untrapped,
};

inline constexpr auto trapped = device_exception_mode::trapped;
inline constexpr auto untrapped = device_exception_mode::untrapped;

} // namespace boost::safe_numbers

The enum is shared with the CUDA backend for source-level symmetry.

SYCL has no portable device trap analogous to CUDA’s __trap(), so under SYCL error reporting is always deferred: the device records the error and the kernel runs to completion; the error surfaces on the host at synchronize(). trapped is accepted but degrades to untrapped behavior. Code that must abort on the first device error should treat the exception thrown by synchronize() as fatal.

The device_error_context Class

#include <boost/safe_numbers/device_error_reporting.hpp>

namespace boost::safe_numbers {

class device_error_context
{
public:
    explicit device_error_context(sycl::queue q);
    device_error_context(sycl::queue q, device_exception_mode e);
    ~device_error_context();

    device_error_context(const device_error_context&) = delete;
    device_error_context& operator=(const device_error_context&) = delete;

    void reset();
    void set_device_exception_method(device_exception_mode e);
    void synchronize();
};

} // namespace boost::safe_numbers

The device_error_context manages a SYCL device_global buffer used to capture errors from device code. When a safe_numbers operation detects an error on the device (overflow, underflow, domain error, invalid argument), the error details (file, line, work-item id, expression, and exception type) are written into this buffer by the first failing work-item. The host reads the buffer back during synchronize() and throws the corresponding std::exception.

Only one device_error_context may exist at a time. Constructing a second instance while one is already alive throws std::logic_error. This constraint prevents races on the shared error buffer.

Constructors

explicit device_error_context(sycl::queue q);

Constructs a context bound to q. Clears any stale error state on the device.

device_error_context(sycl::queue q, device_exception_mode e);

As above, additionally recording the (advisory, see the note above) exception mode.

reset

void reset();

Zeroes the device error buffer (host to device) so the context can be reused across kernel launches. Called automatically by the constructors and by synchronize() after reading the error state.

set_device_exception_method

void set_device_exception_method(device_exception_mode e);

Records the exception mode after construction. Advisory under SYCL (reporting is always deferred).

synchronize

void synchronize();

Calls queue::wait() to finish all enqueued kernels, reads the device buffer back to the host, and if an error was captured clears it and throws the matching exception:

Device Error Host Exception

Overflow

std::overflow_error

Underflow

std::underflow_error

Domain error (e.g. division by zero, bounded out of range)

std::domain_error

Invalid argument

std::invalid_argument

Unknown

std::runtime_error

The error state is cleared before throwing, so after catching the exception the same context is immediately reusable; no manual reset() call is needed.

Examples

  • SYCL Device Support: runs safe_numbers arithmetic on a SYCL device and verifies the results against the host.

  • SYCL Error Handling: shows how to use device_error_context to catch a device-side error on the host.

Edit this Page