diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index f6e784c..85f3dd1 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -105,6 +105,68 @@ jobs: cd $GITHUB_WORKSPACE done + sycl-cmake-test: + strategy: + fail-fast: false + + runs-on: ubuntu-latest + + steps: + - name: Intel Apt repository + timeout-minutes: 1 + run: | + wget https://apt.repos.intel.com/intel-gpg-keys/GPG-PUB-KEY-INTEL-SW-PRODUCTS-2023.PUB + sudo apt-key add GPG-PUB-KEY-INTEL-SW-PRODUCTS-2023.PUB + rm GPG-PUB-KEY-INTEL-SW-PRODUCTS-2023.PUB + echo "deb https://apt.repos.intel.com/oneapi all main" | sudo tee /etc/apt/sources.list.d/oneAPI.list + sudo apt-get update + - name: Install Intel oneAPI compilers + timeout-minutes: 5 + run: sudo apt-get install -y intel-oneapi-compiler-fortran intel-oneapi-compiler-dpcpp-cpp + + - name: Setup Intel oneAPI environment + run: | + source /opt/intel/oneapi/setvars.sh + printenv >> $GITHUB_ENV + + - uses: actions/checkout@v5 + + - name: Install Packages + run: | + sudo apt-get install -y cmake make + - name: Setup Boost + run: | + echo GITHUB_REPOSITORY: $GITHUB_REPOSITORY + LIBRARY=${GITHUB_REPOSITORY#*/} + echo LIBRARY: $LIBRARY + echo "LIBRARY=$LIBRARY" >> $GITHUB_ENV + echo GITHUB_BASE_REF: $GITHUB_BASE_REF + echo GITHUB_REF: $GITHUB_REF + REF=${GITHUB_BASE_REF:-$GITHUB_REF} + REF=${REF#refs/heads/} + echo REF: $REF + BOOST_BRANCH=develop && [ "$REF" == "master" ] && BOOST_BRANCH=master || true + echo BOOST_BRANCH: $BOOST_BRANCH + cd .. + git clone -b $BOOST_BRANCH --depth 1 https://github.com/boostorg/boost.git boost-root + cd boost-root + mkdir -p libs/$LIBRARY + cp -r $GITHUB_WORKSPACE/* libs/$LIBRARY + git submodule update --init tools/boostdep + python3 tools/boostdep/depinst/depinst.py --git_args "--jobs 3" $LIBRARY + - name: Test C++20 + run: | + for std in 20; do + echo "======== Testing C++${std} ========" + cd ../boost-root + rm -rf __build__ + mkdir __build__ && cd __build__ + cmake -DBOOST_INCLUDE_LIBRARIES=$LIBRARY -DBUILD_TESTING=ON -DBOOST_SAFE_NUMBERS_ENABLE_SYCL=ON -DCMAKE_CXX_COMPILER=icpx -DCMAKE_C_COMPILER=icx -DCMAKE_CXX_STANDARD=${std} .. + cmake --build . --target tests -j $(nproc) + ctest --output-on-failure --no-tests=error + cd $GITHUB_WORKSPACE + done + compile-time-checks: name: Compile-time checks (GCC) runs-on: ubuntu-latest diff --git a/doc/modules/ROOT/nav.adoc b/doc/modules/ROOT/nav.adoc index a7a366a..37a22a7 100644 --- a/doc/modules/ROOT/nav.adoc +++ b/doc/modules/ROOT/nav.adoc @@ -27,6 +27,8 @@ ** xref:examples.adoc#examples_cuda_error_handling[CUDA Error Handling] *** xref:examples.adoc#examples_cuda_error_handling_with_error_context[With Safe_Numbers Error Context] *** xref:examples.adoc#examples_cuda_error_handling_without_error_context[Without Safe_Numbers Error Context] +** xref:examples.adoc#examples_sycl[SYCL Device Support] +** xref:examples.adoc#examples_sycl_error_handling[SYCL Error Handling] * xref:pretty_printers.adoc[] * xref:api_reference.adoc[] ** xref:api_reference.adoc#api_namespaces[Namespaces] @@ -41,6 +43,7 @@ * xref:bounded_int.adoc[] * xref:bounded_float.adoc[] * xref:cuda.adoc[] +* xref:sycl.adoc[] * xref:literals.adoc[] * Standard Library Support ** xref:limits.adoc[] diff --git a/doc/modules/ROOT/pages/api_reference.adoc b/doc/modules/ROOT/pages/api_reference.adoc index 241aa95..148db89 100644 --- a/doc/modules/ROOT/pages/api_reference.adoc +++ b/doc/modules/ROOT/pages/api_reference.adoc @@ -136,6 +136,35 @@ https://www.boost.org/LICENSE_1_0.txt | Synchronizes the device, checks for captured errors, and throws the corresponding host exception |=== +=== SYCL Support + +[cols="1,2", options="header"] +|=== +| Type | Description + +| xref:sycl.adoc#sycl_device_error_context[`device_error_context`] +| SYCL error context bound to a queue; the device records errors implicitly and `synchronize()` rethrows them on the host + +| xref:sycl.adoc#sycl_device_exception_mode[`device_exception_mode`] +| Shared with the CUDA backend; always deferred under SYCL +|=== + +==== `device_error_context` Member Functions (SYCL) + +[cols="1,2", options="header"] +|=== +| Function | Description + +| xref:sycl.adoc#sycl_device_error_context_reset[`reset`] +| Zeroes the device error buffer so the context can be reused across launches + +| xref:sycl.adoc#sycl_device_error_context_set_mode[`set_device_exception_method`] +| Records the exception mode after construction (advisory under SYCL) + +| xref:sycl.adoc#sycl_device_error_context_synchronize[`synchronize`] +| Waits for the queue, reads the device buffer, and throws the matching host exception +|=== + [#api_functions] == Functions @@ -465,6 +494,6 @@ This header is not included in the convenience header since it requires external | `` | Byte order conversion functions (`to_be`, `from_be`, `to_le`, `from_le`, `to_be_bytes`, `from_be_bytes`, `to_le_bytes`, `from_le_bytes`, `to_ne_bytes`, `from_ne_bytes`) -| `` -| CUDA device error handling (`device_exception_mode`, `device_error_context`) +| `` +| CUDA and SYCL device error handling (`device_exception_mode`, `device_error_context`) |=== diff --git a/doc/modules/ROOT/pages/cuda.adoc b/doc/modules/ROOT/pages/cuda.adoc index edf3579..d8a279b 100644 --- a/doc/modules/ROOT/pages/cuda.adoc +++ b/doc/modules/ROOT/pages/cuda.adoc @@ -145,7 +145,7 @@ See xref:examples.adoc#examples_cuda_bounded[the bounded-types CUDA example]. [source,c++] ---- -#include +#include namespace boost::safe_numbers { @@ -191,7 +191,7 @@ boost::safe_numbers::device_error_context ctx{boost::safe_numbers::untrapped}; [source,c++] ---- -#include +#include namespace boost::safe_numbers { diff --git a/doc/modules/ROOT/pages/examples.adoc b/doc/modules/ROOT/pages/examples.adoc index 8aeea6a..c177429 100644 --- a/doc/modules/ROOT/pages/examples.adoc +++ b/doc/modules/ROOT/pages/examples.adoc @@ -848,3 +848,45 @@ Kernel has been reset via CUDA API Kernel failed with error: CUDA-capable device(s) is/are busy or unavailable ---- ==== + +[#examples_sycl] +== SYCL Device Support + +With `BOOST_SAFE_NUMBERS_ENABLE_SYCL` defined and `` included first, the safe_numbers types run inside SYCL device kernels compiled with `icpx -fsycl`. +This example runs `u128` arithmetic element-wise on a SYCL device over shared memory and verifies the device results against a host recomputation. + +.This https://github.com/boostorg/safe_numbers/blob/develop/examples/sycl.cpp[example] demonstrates safe_numbers arithmetic running on a SYCL device. +==== +[source, c++] +---- +include::example$sycl.cpp[] +---- + +Output: +---- +Running on: Intel(R) Data Center GPU +All CPU and GPU computed elements match! +---- +==== + +[#examples_sycl_error_handling] +== SYCL Error Handling + +A SYCL device kernel cannot throw, so a safe_numbers error is recorded on the device and rethrown on the host by `device_error_context::synchronize()`. +This example forces a `bounded_uint` result out of its declared range on the device and catches the resulting `std::domain_error` on the host. + +.This https://github.com/boostorg/safe_numbers/blob/develop/examples/sycl_error_handling.cpp[example] demonstrates catching a device-side error on the host using `device_error_context`. +==== +[source, c++] +---- +include::example$sycl_error_handling.cpp[] +---- + +Output: +---- +Running on: Intel(R) Data Center GPU +Caught expected device error on the host: Device error on thread 0 at .../boost/safe_numbers/include/boost/safe_numbers/bounded_integers.hpp:215: bounded_uint addition result out of range +---- +==== + +The reported thread is the failing work-item's global linear id (here `0`, since every element fails and the first work-item wins the recording race); with a single bad input element the id pinpoints it exactly, as on CUDA. diff --git a/doc/modules/ROOT/pages/sycl.adoc b/doc/modules/ROOT/pages/sycl.adoc new file mode 100644 index 0000000..4d9b6d5 --- /dev/null +++ b/doc/modules/ROOT/pages/sycl.adoc @@ -0,0 +1,238 @@ +//// +Copyright 2026 Matt Borland +Distributed under the Boost Software License, Version 1.0. +https://www.boost.org/LICENSE_1_0.txt +//// + +[#sycl] += SYCL Support +:idprefix: sycl_ + +== 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 `` *before* any Boost.safe_numbers header, so that the `SYCL_EXTERNAL` keyword is available: + +[source, c++] +---- +#include // must come first +#include +#include +---- + +Build with `icpx -fsycl -DBOOST_SAFE_NUMBERS_ENABLE_SYCL=1`. + +[NOTE] +==== +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 `-D__SYCL_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: + +[source, c++] +---- +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(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"). + +[NOTE] +==== +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 `` 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. +==== + +[#sycl_floats] +== 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 (`pass:[+]`, `-`, `pass:[*]`, 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 xref:floats.adoc[]. + +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. + +[#sycl_device_exception_mode] +== The `device_exception_mode` Enum + +[source,c++] +---- +#include + +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 xref:cuda.adoc[CUDA backend] for source-level symmetry. + +[NOTE] +==== +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. +==== + +[#sycl_device_error_context] +== The `device_error_context` Class + +[source,c++] +---- +#include + +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 + +[source,c++] +---- +explicit device_error_context(sycl::queue q); +---- + +Constructs a context bound to `q`. Clears any stale error state on the device. + +[source,c++] +---- +device_error_context(sycl::queue q, device_exception_mode e); +---- + +As above, additionally recording the (advisory, see the note above) exception mode. + +[#sycl_device_error_context_reset] +=== `reset` + +[source,c++] +---- +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. + +[#sycl_device_error_context_set_mode] +=== `set_device_exception_method` + +[source,c++] +---- +void set_device_exception_method(device_exception_mode e); +---- + +Records the exception mode after construction. Advisory under SYCL (reporting is always deferred). + +[#sycl_device_error_context_synchronize] +=== `synchronize` + +[source,c++] +---- +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 + +* xref:examples.adoc#examples_sycl[SYCL Device Support]: runs safe_numbers arithmetic on a SYCL device and verifies the results against the host. +* xref:examples.adoc#examples_sycl_error_handling[SYCL Error Handling]: shows how to use `device_error_context` to catch a device-side error on the host. diff --git a/examples/cuda.cu b/examples/cuda.cu index f1d2724..9d5be1f 100644 --- a/examples/cuda.cu +++ b/examples/cuda.cu @@ -11,7 +11,7 @@ #include #include #include -#include +#include #include diff --git a/examples/cuda_bounded.cu b/examples/cuda_bounded.cu index 18d136d..774c257 100644 --- a/examples/cuda_bounded.cu +++ b/examples/cuda_bounded.cu @@ -11,7 +11,7 @@ #include #include #include -#include +#include #include diff --git a/examples/cuda_error_handling.cu b/examples/cuda_error_handling.cu index c8aac30..a15fa65 100644 --- a/examples/cuda_error_handling.cu +++ b/examples/cuda_error_handling.cu @@ -16,7 +16,7 @@ #include #include #include -#include +#include #include diff --git a/examples/cuda_error_handling_without_error_context.cu b/examples/cuda_error_handling_without_error_context.cu index eeb704d..2313018 100644 --- a/examples/cuda_error_handling_without_error_context.cu +++ b/examples/cuda_error_handling_without_error_context.cu @@ -16,7 +16,7 @@ #include #include #include -#include +#include #include diff --git a/examples/sycl.cpp b/examples/sycl.cpp new file mode 100644 index 0000000..e7285bf --- /dev/null +++ b/examples/sycl.cpp @@ -0,0 +1,79 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +// Demonstrates running Boost.safe_numbers arithmetic on a SYCL device. +// Build (Intel oneAPI): icpx -fsycl -DBOOST_SAFE_NUMBERS_ENABLE_SYCL=1 sycl.cpp +// +// must be included before any safe_numbers header so that the +// SYCL_EXTERNAL keyword is available. + +#include +#include + +#include +#include +#include +#include + +using boost::safe_numbers::u128; + +int main() +{ + sycl::queue q; + std::cout << "Running on: " + << q.get_device().get_info() << "\n"; + + constexpr int num_elements {50000}; + + u128* in1 {sycl::malloc_shared(num_elements, q)}; + u128* in2 {sycl::malloc_shared(num_elements, q)}; + u128* out {sycl::malloc_shared(num_elements, q)}; + + // Inputs kept well below the 128-bit maximum so 2*a + b cannot overflow. + std::mt19937_64 rng {42}; + for (int i {0}; i < num_elements; ++i) + { + in1[i] = u128{rng()}; + in2[i] = u128{rng()}; + } + + q.submit([&](sycl::handler& h) + { + h.parallel_for(sycl::range<1>(num_elements), [=](sycl::id<1> idx) + { + const int i {static_cast(idx[0])}; + out[i] = in1[i] * u128{2} + in2[i]; + }); + }).wait(); + + std::vector expected; + expected.reserve(num_elements); + for (int i {0}; i < num_elements; ++i) + { + expected.push_back(in1[i] * u128{2} + in2[i]); + } + + bool all_match {true}; + for (int i {0}; i < num_elements; ++i) + { + if (out[i] != expected[i]) + { + all_match = false; + break; + } + } + + sycl::free(in1, q); + sycl::free(in2, q); + sycl::free(out, q); + + if (all_match) + { + std::cout << "All CPU and GPU computed elements match!\n"; + return EXIT_SUCCESS; + } + + std::cout << "Mismatch between CPU and GPU results\n"; + return EXIT_FAILURE; +} diff --git a/examples/sycl_error_handling.cpp b/examples/sycl_error_handling.cpp new file mode 100644 index 0000000..5562176 --- /dev/null +++ b/examples/sycl_error_handling.cpp @@ -0,0 +1,65 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +// Demonstrates Boost.safe_numbers device error handling under SYCL. +// A safe_numbers operation cannot throw on a device, so an error (here a +// bounded value leaving its range) is recorded on the device and re-raised on +// the host as a normal std::exception by device_error_context::synchronize(). +// Build (Intel oneAPI): icpx -fsycl -DBOOST_SAFE_NUMBERS_ENABLE_SYCL=1 sycl_error_handling.cpp + +#include +#include +#include + +#include +#include + +using bu = boost::safe_numbers::bounded_uint<0u, 1000u>; + +int main() +{ + sycl::queue q; + std::cout << "Running on: " + << q.get_device().get_info() << "\n"; + + constexpr int num_elements {1024}; + bu* a {sycl::malloc_shared(num_elements, q)}; + bu* b {sycl::malloc_shared(num_elements, q)}; + bu* out {sycl::malloc_shared(num_elements, q)}; + + // 600 + 600 = 1200, which is outside the declared range [0, 1000]. + for (int i {0}; i < num_elements; ++i) + { + a[i] = bu{600}; + b[i] = bu{600}; + } + + // The context must be constructed around the launch and only one may exist + // at a time. It binds to the queue used for the kernels. + 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(idx[0])}; + out[i] = a[i] + b[i]; + }); + + int result {EXIT_FAILURE}; + try + { + // Waits for the kernels, then throws if the device recorded an error. + ctx.synchronize(); + std::cout << "No error was reported (unexpected)\n"; + } + catch (const std::domain_error& e) + { + std::cout << "Caught expected device error on the host: " << e.what() << "\n"; + result = EXIT_SUCCESS; + } + + sycl::free(a, q); + sycl::free(b, q); + sycl::free(out, q); + return result; +} diff --git a/include/boost/safe_numbers/bounded_floats.hpp b/include/boost/safe_numbers/bounded_floats.hpp index 374235c..cb7cf71 100644 --- a/include/boost/safe_numbers/bounded_floats.hpp +++ b/include/boost/safe_numbers/bounded_floats.hpp @@ -61,7 +61,7 @@ class bounded_float // NaN comparisons are unordered: a naked range check would silently accept NaN. if (detail::impl::constexpr_isnan(raw)) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_float NaN value"); // LCOV_EXCL_LINE @@ -78,7 +78,7 @@ class bounded_float if (raw < min_raw || raw > max_raw) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_float value out of range"); // LCOV_EXCL_LINE @@ -107,7 +107,7 @@ class bounded_float if (detail::impl::constexpr_isinf(result) && !detail::impl::constexpr_isinf(raw)) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::overflow_error("bounded_float narrowing conversion overflow"); // LCOV_EXCL_LINE diff --git a/include/boost/safe_numbers/bounded_integers.hpp b/include/boost/safe_numbers/bounded_integers.hpp index 4b0db49..29b0167 100644 --- a/include/boost/safe_numbers/bounded_integers.hpp +++ b/include/boost/safe_numbers/bounded_integers.hpp @@ -95,7 +95,7 @@ class bounded_uint if (val < min_val || val > max_val) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint value out of range"); // LCOV_EXCL_LINE @@ -123,7 +123,7 @@ class bounded_uint using raw_other = detail::underlying_type_t; if (raw > static_cast>(std::numeric_limits::max())) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint conversion overflow"); // LCOV_EXCL_LINE @@ -190,7 +190,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator+(const boun underlying res {}; if (detail::impl::unsigned_no_intrin_add(lhs_raw, rhs_raw, res)) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::overflow_error("bounded_uint addition overflow"); // LCOV_EXCL_LINE @@ -204,7 +204,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator+(const boun if (detail::out_of_bounds(res, min_raw, max_raw)) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint addition result out of range"); // LCOV_EXCL_LINE @@ -233,7 +233,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator-(const boun underlying res {}; if (detail::impl::unsigned_no_intrin_sub(lhs_raw, rhs_raw, res)) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::underflow_error("bounded_uint subtraction underflow"); // LCOV_EXCL_LINE @@ -247,7 +247,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator-(const boun if (detail::out_of_bounds(res, min_raw, max_raw)) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint subtraction result out of range"); // LCOV_EXCL_LINE @@ -276,7 +276,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator*(const boun underlying res {}; if (detail::impl::no_intrin_mul(lhs_raw, rhs_raw, res)) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::overflow_error("bounded_uint multiplication overflow"); // LCOV_EXCL_LINE @@ -290,7 +290,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator*(const boun if (detail::out_of_bounds(res, min_raw, max_raw)) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint multiplication result out of range"); // LCOV_EXCL_LINE @@ -318,7 +318,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator/(const boun if (rhs_raw == 0U) [[unlikely]] { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint division by zero"); // LCOV_EXCL_LINE @@ -342,7 +342,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator/(const boun if (detail::out_of_bounds(res, min_raw, max_raw)) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint division result out of range"); // LCOV_EXCL_LINE @@ -370,7 +370,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator%(const boun if (rhs_raw == 0U) [[unlikely]] { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint modulo by zero"); // LCOV_EXCL_LINE @@ -394,7 +394,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator%(const boun if (detail::out_of_bounds(res, min_raw, max_raw)) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint modulo result out of range"); // LCOV_EXCL_LINE @@ -462,7 +462,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_uint::operator++ underlying res {}; if (detail::impl::unsigned_no_intrin_add(raw, static_cast(1U), res)) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::overflow_error("bounded_uint increment overflow"); // LCOV_EXCL_LINE @@ -476,7 +476,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_uint::operator++ if (res > max_raw) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint increment result out of range"); // LCOV_EXCL_LINE @@ -516,7 +516,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_uint::operator-- underlying res {}; if (detail::impl::unsigned_no_intrin_sub(raw, static_cast(1U), res)) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::underflow_error("bounded_uint decrement underflow"); // LCOV_EXCL_LINE @@ -530,7 +530,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_uint::operator-- if (detail::below_bound(res, min_raw)) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint decrement result out of range"); // LCOV_EXCL_LINE @@ -602,7 +602,7 @@ class bounded_int if (val < min_val || val > max_val) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int value out of range"); // LCOV_EXCL_LINE @@ -631,7 +631,7 @@ class bounded_int if (raw > static_cast>(std::numeric_limits::max()) || raw < static_cast>(std::numeric_limits::min())) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int conversion overflow"); // LCOV_EXCL_LINE @@ -703,7 +703,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_int::operator-() if (raw == std::numeric_limits::min()) [[unlikely]] { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::overflow_error("bounded_int negation overflow"); // LCOV_EXCL_LINE @@ -738,7 +738,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator+(const boun const auto status {detail::impl::signed_no_intrin_add(lhs_raw, rhs_raw, res)}; if (status == detail::impl::signed_overflow_status::overflow) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::overflow_error("bounded_int addition overflow"); // LCOV_EXCL_LINE @@ -751,7 +751,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator+(const boun } else if (status == detail::impl::signed_overflow_status::underflow) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::underflow_error("bounded_int addition underflow"); // LCOV_EXCL_LINE @@ -765,7 +765,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator+(const boun if (detail::out_of_bounds(res, min_raw, max_raw)) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int addition result out of range"); // LCOV_EXCL_LINE @@ -799,7 +799,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator-(const boun const auto status {detail::impl::signed_no_intrin_sub(lhs_raw, rhs_raw, res)}; if (status == detail::impl::signed_overflow_status::overflow) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::overflow_error("bounded_int subtraction overflow"); // LCOV_EXCL_LINE @@ -812,7 +812,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator-(const boun } else if (status == detail::impl::signed_overflow_status::underflow) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::underflow_error("bounded_int subtraction underflow"); // LCOV_EXCL_LINE @@ -826,7 +826,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator-(const boun if (detail::out_of_bounds(res, min_raw, max_raw)) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int subtraction result out of range"); // LCOV_EXCL_LINE @@ -860,7 +860,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator*(const boun const auto status {detail::impl::signed_no_intrin_mul(lhs_raw, rhs_raw, res)}; if (status == detail::impl::signed_overflow_status::overflow) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::overflow_error("bounded_int multiplication overflow"); // LCOV_EXCL_LINE @@ -873,7 +873,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator*(const boun } else if (status == detail::impl::signed_overflow_status::underflow) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::underflow_error("bounded_int multiplication underflow"); // LCOV_EXCL_LINE @@ -887,7 +887,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator*(const boun if (detail::out_of_bounds(res, min_raw, max_raw)) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int multiplication result out of range"); // LCOV_EXCL_LINE @@ -919,7 +919,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator/(const boun if (rhs_raw == static_cast(0)) [[unlikely]] { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int division by zero"); // LCOV_EXCL_LINE @@ -934,7 +934,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator/(const boun if (lhs_raw == std::numeric_limits::min() && rhs_raw == static_cast(-1)) [[unlikely]] { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::overflow_error("bounded_int division overflow"); // LCOV_EXCL_LINE @@ -964,7 +964,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator/(const boun if (detail::out_of_bounds(res, min_raw, max_raw)) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int division result out of range"); // LCOV_EXCL_LINE @@ -996,7 +996,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator%(const boun if (rhs_raw == static_cast(0)) [[unlikely]] { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int modulo by zero"); // LCOV_EXCL_LINE @@ -1011,7 +1011,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator%(const boun if (lhs_raw == std::numeric_limits::min() && rhs_raw == static_cast(-1)) [[unlikely]] { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::overflow_error("bounded_int modulo overflow"); // LCOV_EXCL_LINE @@ -1041,7 +1041,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator%(const boun if (detail::out_of_bounds(res, min_raw, max_raw)) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int modulo result out of range"); // LCOV_EXCL_LINE @@ -1118,7 +1118,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_int::operator++( const auto status {detail::impl::signed_no_intrin_add(raw, static_cast(1), res)}; if (status != detail::impl::signed_overflow_status::no_error) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::overflow_error("bounded_int increment overflow"); // LCOV_EXCL_LINE @@ -1132,7 +1132,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_int::operator++( if (res > max_raw) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int increment result out of range"); // LCOV_EXCL_LINE @@ -1173,7 +1173,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_int::operator--( const auto status {detail::impl::signed_no_intrin_sub(raw, static_cast(1), res)}; if (status != detail::impl::signed_overflow_status::no_error) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::underflow_error("bounded_int decrement underflow"); // LCOV_EXCL_LINE @@ -1187,7 +1187,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_int::operator--( if (detail::below_bound(res, min_raw)) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int decrement result out of range"); // LCOV_EXCL_LINE diff --git a/include/boost/safe_numbers/cuda_error_reporting.hpp b/include/boost/safe_numbers/cuda_error_reporting.hpp deleted file mode 100644 index eaf8599..0000000 --- a/include/boost/safe_numbers/cuda_error_reporting.hpp +++ /dev/null @@ -1,310 +0,0 @@ -// Copyright 2026 Matt Borland -// Distributed under the Boost Software License, Version 1.0. -// https://www.boost.org/LICENSE_1_0.txt - -#ifndef BOOST_SAFE_NUMBERS_CUDA_ERROR_REPORTING_HPP -#define BOOST_SAFE_NUMBERS_CUDA_ERROR_REPORTING_HPP - -#include - -#ifndef BOOST_SAFE_NUMBERS_BUILD_MODULE - -#include -#include -#include -#include - -#ifdef __CUDACC__ -#include -#endif - -#endif // BOOST_SAFE_NUMBERS_BUILD_MODULE - -// Using a macro instead of a global constant because the inline constexpr is not available on device -#ifdef PATH_MAX -# define BOOST_SAFE_NUMBERS_DEVICE_ERROR_BUFFER_SIZE PATH_MAX -#else -# define BOOST_SAFE_NUMBERS_DEVICE_ERROR_BUFFER_SIZE 512 -#endif - -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 detail { - -enum class exception_type : unsigned -{ - domain_error, - overflow, - underflow, - invalid_argument, - unknown, -}; - -struct cuda_device_error -{ - int flag; // 0 = no error, 1 = error captured - int line; // __LINE__ - int thread_id; // blockIdx.x * blockDim.x + threadIdx.x - exception_type exception; // Type of exception that would have been thrown on CPU - char file[BOOST_SAFE_NUMBERS_DEVICE_ERROR_BUFFER_SIZE]; // __FILE__ copied by value - char expression[BOOST_SAFE_NUMBERS_DEVICE_ERROR_BUFFER_SIZE]; // x copied by value -}; - -// Compile-time map from std exception type to our enum -template -BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto to_exception_enum() noexcept -> exception_type -{ - if constexpr (std::is_same_v) - { - return exception_type::domain_error; - } - else if constexpr (std::is_same_v) - { - return exception_type::overflow; - } - else if constexpr (std::is_same_v) - { - return exception_type::underflow; - } - else if constexpr (std::is_same_v) - { - return exception_type::invalid_argument; - } - else - { - return exception_type::unknown; - } -} - -BOOST_SAFE_NUMBERS_HOST_DEVICE inline void copy_to_buf(char* dst, const char* src, const int max_len) -{ - int i = 0; - for (; i < max_len - 1 && src[i] != '\0'; ++i) - { - dst[i] = src[i]; - } - dst[i] = '\0'; -} - -#ifdef __CUDACC__ - -// Managed memory error struct accessible from both host and device. -// Since we never destroy the CUDA context, __managed__ is safe to use. -__managed__ cuda_device_error g_device_error {}; - -// Managed memory enum class that allows us to set what report_device_error should do -// We default to trapped as that's the best way to ensure hard failure in the event of error -__managed__ device_exception_mode g_device_fail_type {device_exception_mode::trapped}; - -// Tracks whether a device_error_context instance is alive. -// Only one may exist at a time to prevent races on g_device_error. -inline bool g_device_error_context_active = false; - -__host__ __device__ inline void report_device_error( - exception_type exc, - const char* file, - int line, - const char* expression) -{ - #ifdef __CUDA_ARCH__ - - if (atomicCAS(&g_device_error.flag, 0, 1) == 0) - { - g_device_error.line = line; - g_device_error.thread_id = blockIdx.x * blockDim.x + threadIdx.x; - g_device_error.exception = exc; - - copy_to_buf(g_device_error.file, file, BOOST_SAFE_NUMBERS_DEVICE_ERROR_BUFFER_SIZE); - copy_to_buf(g_device_error.expression, expression, BOOST_SAFE_NUMBERS_DEVICE_ERROR_BUFFER_SIZE); - __threadfence_system(); - - if (g_device_fail_type == device_exception_mode::trapped) - { - __trap(); - } - } - - switch (g_device_fail_type) - { - case device_exception_mode::trapped: - // In the event that __trap() is called the error is non-recoverable - // The user must terminate the current PROCESS in order to reuse the device - // There is currently (3/26) way to recover using the cuda_runtime or hardware APIs - // Other threads: spin until the trap terminates the kernel - while (true) - { - __nanosleep(1000000); - } - break; - - case device_exception_mode::untrapped: - // Return instead of calling __trap(). This allows the kernel to - // complete normally without corrupting the CUDA context. Other - // threads may continue with incorrect values, but synchronize() - // will detect the error via the flag and throw on the host. - return; - break; - } - - #else - - const auto msg = std::string(file) + ":" + std::to_string(line) + ": " + expression; - switch (exc) - { - case exception_type::domain_error: - BOOST_THROW_EXCEPTION(std::domain_error(msg)); - break; - case exception_type::overflow: - BOOST_THROW_EXCEPTION(std::overflow_error(msg)); - break; - case exception_type::underflow: - BOOST_THROW_EXCEPTION(std::underflow_error(msg)); - break; - case exception_type::invalid_argument: - BOOST_THROW_EXCEPTION(std::invalid_argument(msg)); - break; - case exception_type::unknown: - [[fallthrough]]; - default: - BOOST_THROW_EXCEPTION(std::runtime_error(msg)); - } - - #endif -} - -#endif // __CUDACC__ - -} // namespace detail - -#ifdef __CUDACC__ - -class device_error_context -{ -public: - - // Clears the error state. Only one device_error_context may exist at a time. - device_error_context() - { - if (detail::g_device_error_context_active) - { - BOOST_THROW_EXCEPTION(std::logic_error( - "Only one device_error_context may exist at a time")); - } - detail::g_device_error_context_active = true; - reset(); - } - - // Sets a different error type to our managed global variable - device_error_context(const device_exception_mode e) - { - if (detail::g_device_error_context_active) - { - BOOST_THROW_EXCEPTION(std::logic_error( - "Only one device_error_context may exist at a time")); - } - - detail::g_device_fail_type = e; - - detail::g_device_error_context_active = true; - reset(); - } - - ~device_error_context() - { - detail::g_device_error_context_active = false; - } - - device_error_context(const device_error_context&) = delete; - device_error_context& operator=(const device_error_context&) = delete; - - // Clears the error fields so the context can be reused across kernel launches. - void reset() - { - detail::g_device_error.flag = 0; - detail::g_device_error.line = 0; - detail::g_device_error.thread_id = 0; - detail::g_device_error.exception = detail::exception_type::unknown; - detail::g_device_error.file[0] = '\0'; - detail::g_device_error.expression[0] = '\0'; - } - - // Adds a post-construction way of setting the failure mode for the device - void set_device_exception_method(const device_exception_mode e) - { - detail::g_device_fail_type = e; - } - - // Synchronizes the device and checks for errors captured by device code. - // If an error was detected, the error state is cleared (so the context - // is immediately reusable), and the appropriate std::exception is thrown. - void synchronize() - { - const auto status = cudaDeviceSynchronize(); - - const auto flag = detail::g_device_error.flag; - - if (flag != 0) - { - const auto thread_id = detail::g_device_error.thread_id; - const auto line = detail::g_device_error.line; - const auto exc = detail::g_device_error.exception; - - std::ostringstream oss; - oss << "Device error on thread " << thread_id - << " at " << detail::g_device_error.file - << ":" << line - << ": " << detail::g_device_error.expression; - - const auto msg = oss.str(); - - // Clear the error state so the context can be reused - // immediately after catching the exception. - reset(); - - switch (exc) - { - case detail::exception_type::domain_error: - BOOST_THROW_EXCEPTION(std::domain_error(msg)); - break; - case detail::exception_type::overflow: - BOOST_THROW_EXCEPTION(std::overflow_error(msg)); - break; - case detail::exception_type::underflow: - BOOST_THROW_EXCEPTION(std::underflow_error(msg)); - break; - case detail::exception_type::invalid_argument: - BOOST_THROW_EXCEPTION(std::invalid_argument(msg)); - break; - case detail::exception_type::unknown: - [[fallthrough]]; - default: - BOOST_THROW_EXCEPTION(std::runtime_error(msg)); - } - } - else - { - reset(); - } - - if (status != cudaSuccess) - { - cudaGetLastError(); - BOOST_THROW_EXCEPTION(std::runtime_error(cudaGetErrorString(status))); - } - } -}; - -#endif // __CUDACC__ - -} // namespace boost::safe_numbers - -#endif // BOOST_SAFE_NUMBERS_CUDA_ERROR_REPORTING_HPP diff --git a/include/boost/safe_numbers/detail/config.hpp b/include/boost/safe_numbers/detail/config.hpp index cf03d48..140fbf0 100644 --- a/include/boost/safe_numbers/detail/config.hpp +++ b/include/boost/safe_numbers/detail/config.hpp @@ -17,7 +17,9 @@ # define BOOST_SAFE_NUMBERS_HAS_BUILTIN(x) 0 #endif // __has_builtin -#ifdef __x86_64__ +// __x86_64__ remains defined on the SYCL spir64 device pass, but these x86 intrinsic +// headers do not compile for that target, so exclude them when __SYCL_DEVICE_ONLY__ is set. +#if defined(__x86_64__) && !defined(__SYCL_DEVICE_ONLY__) #ifndef BOOST_SAFE_NUMBERS_BUILD_MODULE # include @@ -76,6 +78,8 @@ inline constexpr auto dependent_false {false}; } // namespace boost::safe_numbers::detail +// BOOST_SAFE_NUMBERS_HOST_DEVICE annotates every device-callable function and +// expands to "__host__ __device__" (CUDA), "SYCL_EXTERNAL" (SYCL), or nothing. #if defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA) #ifndef BOOST_SAFE_NUMBERS_DETAIL_INT128_ENABLE_CUDA @@ -87,11 +91,32 @@ inline constexpr auto dependent_false {false}; #endif #define BOOST_SAFE_NUMBERS_HOST_DEVICE __host__ __device__ +#define BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT + +#elif defined(BOOST_SAFE_NUMBERS_ENABLE_SYCL) + +// SYCL is fully opt-in. must be included before any safe_numbers +// header so the SYCL_EXTERNAL keyword exists; icpx defines __SYCL_DEVICE_ONLY__ +// on the spir64 device pass, which forces the portable code paths downstream. +#ifndef SYCL_EXTERNAL +# error "Include before any Boost.safe_numbers header when BOOST_SAFE_NUMBERS_ENABLE_SYCL is defined" +#endif + +#ifndef BOOST_SAFE_NUMBERS_DETAIL_INT128_ENABLE_SYCL +# define BOOST_SAFE_NUMBERS_DETAIL_INT128_ENABLE_SYCL +#endif + +#ifndef BOOST_CHARCONV_ENABLE_SYCL +# define BOOST_CHARCONV_ENABLE_SYCL +#endif + +#define BOOST_SAFE_NUMBERS_HOST_DEVICE SYCL_EXTERNAL +#define BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT #else #define BOOST_SAFE_NUMBERS_HOST_DEVICE -#endif // CUDA +#endif // GPU backends #endif // BOOST_SAFENUMBERS_CONFIG_HPP diff --git a/include/boost/safe_numbers/detail/cuda_error_reporting.hpp b/include/boost/safe_numbers/detail/cuda_error_reporting.hpp new file mode 100644 index 0000000..f76ff8c --- /dev/null +++ b/include/boost/safe_numbers/detail/cuda_error_reporting.hpp @@ -0,0 +1,201 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// CUDA device error-reporting backend. Selected by the public +// device_error_reporting.hpp dispatcher when compiling with nvcc and +// BOOST_SAFE_NUMBERS_ENABLE_CUDA. Backend-neutral pieces live in +// detail/device_error_common.hpp. + +#ifndef BOOST_SAFE_NUMBERS_DETAIL_CUDA_ERROR_REPORTING_HPP +#define BOOST_SAFE_NUMBERS_DETAIL_CUDA_ERROR_REPORTING_HPP + +#include +#include + +#ifndef BOOST_SAFE_NUMBERS_BUILD_MODULE + +#include +#include +#include +#include + +#ifdef __CUDACC__ +#include +#endif + +#endif // BOOST_SAFE_NUMBERS_BUILD_MODULE + +namespace boost::safe_numbers { + +namespace detail { + +#ifdef __CUDACC__ + +// Managed memory error record accessible from both host and device. +// Since we never destroy the CUDA context, __managed__ is safe to use. +__managed__ device_error_record g_device_error {}; + +// Managed enum controlling what report_device_error does on failure. +// Default to trapped as that's the best way to ensure hard failure on error. +__managed__ device_exception_mode g_device_fail_type {device_exception_mode::trapped}; + +__host__ __device__ inline void report_device_error( + exception_type exc, + const char* file, + int line, + const char* expression) +{ + #ifdef __CUDA_ARCH__ + + if (atomicCAS(&g_device_error.flag, 0, 1) == 0) + { + g_device_error.line = line; + g_device_error.thread_id = blockIdx.x * blockDim.x + threadIdx.x; + g_device_error.exception = exc; + + copy_to_buf(g_device_error.file, file, BOOST_SAFE_NUMBERS_DEVICE_ERROR_BUFFER_SIZE); + copy_to_buf(g_device_error.expression, expression, BOOST_SAFE_NUMBERS_DEVICE_ERROR_BUFFER_SIZE); + __threadfence_system(); + + if (g_device_fail_type == device_exception_mode::trapped) + { + __trap(); + } + } + + switch (g_device_fail_type) + { + case device_exception_mode::trapped: + // In the event that __trap() is called the error is non-recoverable. + // Other threads: spin until the trap terminates the kernel. + while (true) + { + __nanosleep(1000000); + } + break; + + case device_exception_mode::untrapped: + // Return instead of trapping. The kernel completes; synchronize() + // detects the error via the flag and throws on the host. + return; + break; + } + + #else + + const auto msg = std::string(file) + ":" + std::to_string(line) + ": " + expression; + throw_from_exception_type(exc, msg); + + #endif +} + +#endif // __CUDACC__ + +} // namespace detail + +#ifdef __CUDACC__ + +class device_error_context +{ +public: + + // Clears the error state. Only one device_error_context may exist at a time. + device_error_context() + { + if (detail::g_device_error_context_active) + { + BOOST_THROW_EXCEPTION(std::logic_error( + "Only one device_error_context may exist at a time")); + } + detail::g_device_error_context_active = true; + reset(); + } + + // Sets the failure mode on the managed global up front. + device_error_context(const device_exception_mode e) + { + if (detail::g_device_error_context_active) + { + BOOST_THROW_EXCEPTION(std::logic_error( + "Only one device_error_context may exist at a time")); + } + + detail::g_device_fail_type = e; + + detail::g_device_error_context_active = true; + reset(); + } + + ~device_error_context() + { + detail::g_device_error_context_active = false; + } + + device_error_context(const device_error_context&) = delete; + device_error_context& operator=(const device_error_context&) = delete; + + // Clears the error fields so the context can be reused across kernel launches. + void reset() + { + detail::g_device_error.flag = 0; + detail::g_device_error.line = 0; + detail::g_device_error.thread_id = 0; + detail::g_device_error.exception = detail::exception_type::unknown; + detail::g_device_error.file[0] = '\0'; + detail::g_device_error.expression[0] = '\0'; + } + + // Post-construction way of setting the failure mode for the device. + void set_device_exception_method(const device_exception_mode e) + { + detail::g_device_fail_type = e; + } + + // Synchronizes the device and checks for errors captured by device code. + // If an error was detected, the error state is cleared (so the context is + // immediately reusable), and the appropriate std::exception is thrown. + void synchronize() + { + const auto status = cudaDeviceSynchronize(); + + const auto flag = detail::g_device_error.flag; + + if (flag != 0) + { + const auto thread_id = detail::g_device_error.thread_id; + const auto line = detail::g_device_error.line; + const auto exc = detail::g_device_error.exception; + + std::ostringstream oss; + oss << "Device error on thread " << thread_id + << " at " << detail::g_device_error.file + << ":" << line + << ": " << detail::g_device_error.expression; + + const auto msg = oss.str(); + + // Clear the error state so the context can be reused + // immediately after catching the exception. + reset(); + + detail::throw_from_exception_type(exc, msg); + } + else + { + reset(); + } + + if (status != cudaSuccess) + { + cudaGetLastError(); + BOOST_THROW_EXCEPTION(std::runtime_error(cudaGetErrorString(status))); + } + } +}; + +#endif // __CUDACC__ + +} // namespace boost::safe_numbers + +#endif // BOOST_SAFE_NUMBERS_DETAIL_CUDA_ERROR_REPORTING_HPP diff --git a/include/boost/safe_numbers/detail/device_error_common.hpp b/include/boost/safe_numbers/detail/device_error_common.hpp new file mode 100644 index 0000000..e911365 --- /dev/null +++ b/include/boost/safe_numbers/detail/device_error_common.hpp @@ -0,0 +1,139 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// Backend-neutral pieces shared by the CUDA and SYCL device error reporters. +// Included by detail/cuda_error_reporting.hpp, detail/sycl_error_reporting.hpp, +// and the public device_error_reporting.hpp dispatcher. + +#ifndef BOOST_SAFE_NUMBERS_DETAIL_DEVICE_ERROR_COMMON_HPP +#define BOOST_SAFE_NUMBERS_DETAIL_DEVICE_ERROR_COMMON_HPP + +#include + +#ifndef BOOST_SAFE_NUMBERS_BUILD_MODULE + +#include +#include +#include +#include + +#endif // BOOST_SAFE_NUMBERS_BUILD_MODULE + +// Using a macro instead of a global constant because inline constexpr is not available on device +#ifdef PATH_MAX +# define BOOST_SAFE_NUMBERS_DEVICE_ERROR_BUFFER_SIZE PATH_MAX +#else +# define BOOST_SAFE_NUMBERS_DEVICE_ERROR_BUFFER_SIZE 512 +#endif + +namespace boost::safe_numbers { + +// Selects how the device reacts to a captured error. On CUDA, trapped hard-fails +// the kernel via __trap() while untrapped defers to the host. SYCL has no +// recoverable device trap, so under SYCL it is always deferred (trapped degrades +// to untrapped); the enum is retained for source-level symmetry with CUDA. +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 detail { + +enum class exception_type : unsigned +{ + domain_error, + overflow, + underflow, + invalid_argument, + unknown, +}; + +// POD error record copied by value between device and host. +struct device_error_record +{ + int flag; // 0 = no error, 1 = error captured + int line; // __LINE__ + int thread_id; // thread / work-item id (-1 = unknown) + exception_type exception; // exception that would have been thrown on CPU + char file[BOOST_SAFE_NUMBERS_DEVICE_ERROR_BUFFER_SIZE]; // __FILE__ copied by value + char expression[BOOST_SAFE_NUMBERS_DEVICE_ERROR_BUFFER_SIZE]; // message copied by value +}; + +// Back-compat alias for the original CUDA-specific name. +using cuda_device_error = device_error_record; + +// Compile-time map from a std exception type to our enum. +template +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto to_exception_enum() noexcept -> exception_type +{ + if constexpr (std::is_same_v) + { + return exception_type::domain_error; + } + else if constexpr (std::is_same_v) + { + return exception_type::overflow; + } + else if constexpr (std::is_same_v) + { + return exception_type::underflow; + } + else if constexpr (std::is_same_v) + { + return exception_type::invalid_argument; + } + else + { + return exception_type::unknown; + } +} + +// Device-safe bounded string copy (device has no strncpy). +BOOST_SAFE_NUMBERS_HOST_DEVICE inline void copy_to_buf(char* dst, const char* src, const int max_len) +{ + int i = 0; + for (; i < max_len - 1 && src[i] != '\0'; ++i) + { + dst[i] = src[i]; + } + dst[i] = '\0'; +} + +// Host-side translation from the enum to a real std exception. Used by both +// backends' host branch and by both device_error_context::synchronize() methods. +inline void throw_from_exception_type(const exception_type exc, const std::string& msg) +{ + switch (exc) + { + case exception_type::domain_error: + BOOST_THROW_EXCEPTION(std::domain_error(msg)); + break; + case exception_type::overflow: + BOOST_THROW_EXCEPTION(std::overflow_error(msg)); + break; + case exception_type::underflow: + BOOST_THROW_EXCEPTION(std::underflow_error(msg)); + break; + case exception_type::invalid_argument: + BOOST_THROW_EXCEPTION(std::invalid_argument(msg)); + break; + case exception_type::unknown: + [[fallthrough]]; + default: + BOOST_THROW_EXCEPTION(std::runtime_error(msg)); + } +} + +// Tracks whether a device_error_context instance is alive. Only one may exist at +// a time to prevent races on the device error buffer. +inline bool g_device_error_context_active = false; + +} // namespace detail +} // namespace boost::safe_numbers + +#endif // BOOST_SAFE_NUMBERS_DETAIL_DEVICE_ERROR_COMMON_HPP diff --git a/include/boost/safe_numbers/detail/float_basis.hpp b/include/boost/safe_numbers/detail/float_basis.hpp index 2a2d399..0c9225d 100644 --- a/include/boost/safe_numbers/detail/float_basis.hpp +++ b/include/boost/safe_numbers/detail/float_basis.hpp @@ -500,7 +500,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto classify_cmath_resul return res > 0 ? error_category::overflow : error_category::underflow; } -#if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) +#if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) // Host-only descriptive runtime message, e.g. "Overflow detected in f64 sqrt". template @@ -514,7 +514,7 @@ template template BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_cmath_overflow(const char* const op) -> void { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -538,7 +538,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_cmath_overflow(const char* c template BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_cmath_underflow(const char* const op) -> void { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -562,7 +562,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_cmath_underflow(const char* template BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_cmath_domain(const char* const op) -> void { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -678,7 +678,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto checked_float_additi template BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_overflow_add() -> void { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -700,7 +700,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_overflow_add() -> void template BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_underflow_add() -> void { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -722,7 +722,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_underflow_add() -> void template BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_nan_add() -> void { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -744,7 +744,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_nan_add() -> void template BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_invalid_add() -> void { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -781,7 +781,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE case impl::error_category::no_error: break; case impl::error_category::overflow: - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -800,7 +800,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE } break; case impl::error_category::underflow: - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -819,7 +819,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE } break; case impl::error_category::nan_op: - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -838,7 +838,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE } break; case impl::error_category::invalid_op: - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -926,7 +926,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto checked_float_subtra template BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_overflow_sub() -> void { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -948,7 +948,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_overflow_sub() -> void template BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_underflow_sub() -> void { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -970,7 +970,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_underflow_sub() -> void template BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_nan_sub() -> void { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -992,7 +992,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_nan_sub() -> void template BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_invalid_sub() -> void { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1029,7 +1029,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE case impl::error_category::no_error: break; case impl::error_category::overflow: - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1048,7 +1048,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE } break; case impl::error_category::underflow: - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1067,7 +1067,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE } break; case impl::error_category::nan_op: - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1086,7 +1086,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE } break; case impl::error_category::invalid_op: - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1184,7 +1184,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto checked_float_multip template BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_overflow_mul() -> void { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1206,7 +1206,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_overflow_mul() -> void template BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_underflow_mul() -> void { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1228,7 +1228,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_underflow_mul() -> void template BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_nan_mul() -> void { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1250,7 +1250,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_nan_mul() -> void template BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_invalid_mul() -> void { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1287,7 +1287,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE case impl::error_category::no_error: break; case impl::error_category::overflow: - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1306,7 +1306,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE } break; case impl::error_category::underflow: - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1325,7 +1325,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE } break; case impl::error_category::nan_op: - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1344,7 +1344,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE } break; case impl::error_category::invalid_op: - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1458,7 +1458,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto checked_float_divisi template BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_overflow_div() -> void { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1480,7 +1480,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_overflow_div() -> void template BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_underflow_div() -> void { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1502,7 +1502,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_underflow_div() -> void template BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_nan_div() -> void { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1524,7 +1524,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_nan_div() -> void template BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_invalid_div() -> void { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1546,7 +1546,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_invalid_div() -> void template BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto throw_divbyzero_div() -> void { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1583,7 +1583,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE case impl::error_category::no_error: break; case impl::error_category::overflow: - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1602,7 +1602,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE } break; case impl::error_category::underflow: - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1621,7 +1621,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE } break; case impl::error_category::nan_op: - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1640,7 +1640,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE } break; case impl::error_category::invalid_op: - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1659,7 +1659,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE } break; case impl::error_category::divide_by_zero: - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) diff --git a/include/boost/safe_numbers/detail/signed_integer_basis.hpp b/include/boost/safe_numbers/detail/signed_integer_basis.hpp index bfe36cc..3420994 100644 --- a/include/boost/safe_numbers/detail/signed_integer_basis.hpp +++ b/include/boost/safe_numbers/detail/signed_integer_basis.hpp @@ -505,7 +505,7 @@ struct signed_add_helper auto handle_error = [&result](signed_overflow_status status) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if (status == signed_overflow_status::overflow) @@ -602,7 +602,7 @@ struct signed_add_helper if constexpr (!std::is_same_v) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (!std::is_constant_evaluated()) { @@ -665,7 +665,7 @@ struct signed_add_helper if constexpr (!std::is_same_v) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (!std::is_constant_evaluated()) { @@ -702,7 +702,7 @@ struct signed_add_helper if constexpr (!std::is_same_v) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (!std::is_constant_evaluated()) { @@ -758,7 +758,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator+(const signed_integer_basis lhs, const signed_integer_basis rhs) -> signed_integer_basis { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { @@ -1137,7 +1137,7 @@ struct signed_sub_helper auto handle_error = [&result](signed_overflow_status status) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if (status == signed_overflow_status::overflow) @@ -1234,7 +1234,7 @@ struct signed_sub_helper if constexpr (!std::is_same_v) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (!std::is_constant_evaluated()) { @@ -1297,7 +1297,7 @@ struct signed_sub_helper if constexpr (!std::is_same_v) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (!std::is_constant_evaluated()) { @@ -1334,7 +1334,7 @@ struct signed_sub_helper if constexpr (!std::is_same_v) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (!std::is_constant_evaluated()) { @@ -1373,7 +1373,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator-(const signed_integer_basis lhs, const signed_integer_basis rhs) -> signed_integer_basis { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { @@ -1737,7 +1737,7 @@ struct signed_mul_helper auto handle_error = [&result](signed_overflow_status status) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if (status == signed_overflow_status::overflow) @@ -1838,7 +1838,7 @@ struct signed_mul_helper if constexpr (!std::is_same_v) #endif { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (!std::is_constant_evaluated()) { @@ -1905,7 +1905,7 @@ struct signed_mul_helper if constexpr (!std::is_same_v) #endif { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (!std::is_constant_evaluated()) { @@ -1946,7 +1946,7 @@ struct signed_mul_helper if constexpr (!std::is_same_v) #endif { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (!std::is_constant_evaluated()) { @@ -2002,7 +2002,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator*(const signed_integer_basis lhs, const signed_integer_basis rhs) -> signed_integer_basis { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { @@ -2165,7 +2165,7 @@ struct signed_div_helper } else { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -2217,7 +2217,7 @@ struct signed_div_helper } else { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -2448,7 +2448,7 @@ struct signed_mod_helper } else { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -2503,7 +2503,7 @@ struct signed_mod_helper } else { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) diff --git a/include/boost/safe_numbers/detail/sycl_error_reporting.hpp b/include/boost/safe_numbers/detail/sycl_error_reporting.hpp new file mode 100644 index 0000000..c288c71 --- /dev/null +++ b/include/boost/safe_numbers/detail/sycl_error_reporting.hpp @@ -0,0 +1,207 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// SYCL device error-reporting backend. Selected by the public +// device_error_reporting.hpp dispatcher when BOOST_SAFE_NUMBERS_ENABLE_SYCL is +// set. Requires to have been included first. Backend-neutral +// pieces live in detail/device_error_common.hpp. + +#ifndef BOOST_SAFE_NUMBERS_DETAIL_SYCL_ERROR_REPORTING_HPP +#define BOOST_SAFE_NUMBERS_DETAIL_SYCL_ERROR_REPORTING_HPP + +#include +#include + +#if defined(BOOST_SAFE_NUMBERS_ENABLE_SYCL) + +#ifndef SYCL_EXTERNAL +# error "Include before when BOOST_SAFE_NUMBERS_ENABLE_SYCL is defined" +#endif + +// The device error buffer (g_device_error below) is a device_global, which core +// SYCL 2020 does not provide - section 5.4 forbids mutable variables with static +// storage duration in device code. device_global is the sycl_ext_oneapi_device_global +// extension that lifts that restriction, and it is the only way a deeply-nested +// operation can reach the shared error record without threading a pointer through +// every call. Require the extension explicitly (its feature-test macro) rather than +// letting device_global fail to resolve with a cryptic diagnostic on implementations +// that do not support it. +#ifndef SYCL_EXT_ONEAPI_DEVICE_GLOBAL +# error "Boost.safe_numbers SYCL device error reporting requires the device_global extension (feature-test macro SYCL_EXT_ONEAPI_DEVICE_GLOBAL), which this SYCL implementation does not advertise" +#endif + +#ifndef BOOST_SAFE_NUMBERS_BUILD_MODULE + +#include +#include +#include +#include +#include + +#endif // BOOST_SAFE_NUMBERS_BUILD_MODULE + +namespace boost::safe_numbers { + +namespace detail { + +// SYCL analog of the CUDA __managed__ error buffer. A single inline instance +// (the SYCL 2020 device_global extension explicitly permits inline) that any +// device-side function can write by name, giving the same zero-pointer-plumbing +// property as __managed__. No device_image_scope property is set, so host +// queue.copy access is unrestricted (the variable is backed by a shared buffer). +inline sycl::ext::oneapi::experimental::device_global g_device_error; + +// Reports an error from device or host. Compiled for both passes of the +// single-source TU: the device pass records into the device_global (first writer +// wins, then returns, since SYCL has no recoverable device trap); the host pass +// throws the matching std exception. Mirrors how the CUDA reporter branches on +// __CUDA_ARCH__. +BOOST_SAFE_NUMBERS_HOST_DEVICE inline void report_device_error( + exception_type exc, + const char* file, + int line, + const char* expression) +{ + #ifdef __SYCL_DEVICE_ONLY__ + + auto& rec = g_device_error.get(); + sycl::atomic_ref flag_ref(rec.flag); + + // First failing work-item wins and records the details; losers never touch + // the payload, so the host can never observe a torn record. + int expected {0}; + if (flag_ref.compare_exchange_strong(expected, 1)) + { + rec.line = line; + rec.exception = exc; + #if defined(SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES) && !defined(BOOST_SAFE_NUMBERS_SYCL_DISABLE_WORKITEM_ID) + // The failing work-item's dimension-0 global linear id, the analog of + // the CUDA backend's blockIdx.x * blockDim.x + threadIdx.x. Core SYCL + // 2020 only passes the id into the kernel lambda, so reaching it from + // here needs the free-function queries extension; on DPC++ it returns + // the correct id for plain range kernels, nd_range kernels, and 0 for + // single_task. Kernels launched over more than one dimension report + // only dimension 0 (same limitation as the CUDA backend). + rec.thread_id = static_cast( + sycl::ext::oneapi::this_work_item::get_nd_item<1>().get_global_linear_id()); + #else + rec.thread_id = -1; + #endif + copy_to_buf(rec.file, file, BOOST_SAFE_NUMBERS_DEVICE_ERROR_BUFFER_SIZE); + copy_to_buf(rec.expression, expression, BOOST_SAFE_NUMBERS_DEVICE_ERROR_BUFFER_SIZE); + } + + // No __trap()/recoverable device abort in SYCL: record and return. + // synchronize() surfaces the error on the host after the kernel completes. + return; + + #else + + const auto msg = std::string(file) + ":" + std::to_string(line) + ": " + expression; + throw_from_exception_type(exc, msg); + + #endif +} + +} // namespace detail + +// Host-facing RAII context placed around SYCL kernel launches. Mirrors the CUDA +// device_error_context but binds to the user's queue (SYCL has no implicit +// device). Only one may exist at a time. Defined on both compiler passes (its +// methods are host-only and never called from device code, so the device pass +// parses but does not codegen them) so that host code referencing it parses on +// the device pass too. +class device_error_context +{ + sycl::queue queue_; + [[maybe_unused]] device_exception_mode mode_ {device_exception_mode::untrapped}; + + static void guard() + { + if (detail::g_device_error_context_active) + { + BOOST_THROW_EXCEPTION(std::logic_error( + "Only one device_error_context may exist at a time")); + } + } + +public: + + explicit device_error_context(sycl::queue q) + : queue_ {std::move(q)} + { + guard(); + detail::g_device_error_context_active = true; + reset(); + } + + // The mode is advisory under SYCL (always deferred); kept for symmetry. + device_error_context(sycl::queue q, const device_exception_mode e) + : queue_ {std::move(q)}, mode_ {e} + { + guard(); + detail::g_device_error_context_active = true; + reset(); + } + + ~device_error_context() + { + detail::g_device_error_context_active = false; + } + + device_error_context(const device_error_context&) = delete; + device_error_context& operator=(const device_error_context&) = delete; + + // Zero the device buffer (host -> device) so the context is reusable. + void reset() + { + const detail::device_error_record zero {}; + queue_.copy(&zero, detail::g_device_error).wait(); + } + + // SYCL is always deferred; retained for source-level symmetry with CUDA. + void set_device_exception_method(const device_exception_mode e) + { + mode_ = e; + } + + // Wait for all enqueued kernels, read the buffer back (device -> host), and + // if an error was captured clear it and throw the matching std::exception. + void synchronize() + { + queue_.wait(); + + detail::device_error_record rec {}; + queue_.copy(detail::g_device_error, &rec).wait(); + + if (rec.flag != 0) + { + std::ostringstream oss; + oss << "Device error on thread " << rec.thread_id + << " at " << rec.file + << ":" << rec.line + << ": " << rec.expression; + + const auto msg = oss.str(); + + // Clear before throwing so the context is reusable. + reset(); + + detail::throw_from_exception_type(rec.exception, msg); + } + else + { + reset(); + } + } +}; + +} // namespace boost::safe_numbers + +#endif // BOOST_SAFE_NUMBERS_ENABLE_SYCL + +#endif // BOOST_SAFE_NUMBERS_DETAIL_SYCL_ERROR_REPORTING_HPP diff --git a/include/boost/safe_numbers/detail/throw_exception.hpp b/include/boost/safe_numbers/detail/throw_exception.hpp index e7e6de0..f44c7b5 100644 --- a/include/boost/safe_numbers/detail/throw_exception.hpp +++ b/include/boost/safe_numbers/detail/throw_exception.hpp @@ -9,7 +9,7 @@ #ifndef BOOST_SAFE_NUMBERS_THROW_EXCEPTION_HPP #define BOOST_SAFE_NUMBERS_THROW_EXCEPTION_HPP -#include +#include #ifndef BOOST_SAFE_NUMBERS_BUILD_MODULE @@ -18,16 +18,17 @@ #endif // BOOST_SAFE_NUMBERS_BUILD_MODULE // Two-argument form: (exception_type, message) -// On host: constructs and throws the exception with the message -// On CUDA device: passes the const char* message to the device error reporter -#ifndef __CUDACC__ +// On host: constructs and throws the exception with the message. +// On a GPU device backend (CUDA or SYCL): forwards to the device error reporter, +// which itself branches host/device internally. +#ifdef BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT -#define BOOST_SAFE_NUMBERS_THROW_EXCEPTION(exc_type, msg) BOOST_THROW_EXCEPTION(exc_type(msg)) +#define BOOST_SAFE_NUMBERS_THROW_EXCEPTION(exc_type, msg) boost::safe_numbers::detail::report_device_error(boost::safe_numbers::detail::to_exception_enum(), __FILE__, __LINE__, msg) #else -#define BOOST_SAFE_NUMBERS_THROW_EXCEPTION(exc_type, msg) boost::safe_numbers::detail::report_device_error(boost::safe_numbers::detail::to_exception_enum(), __FILE__, __LINE__, msg) +#define BOOST_SAFE_NUMBERS_THROW_EXCEPTION(exc_type, msg) BOOST_THROW_EXCEPTION(exc_type(msg)) -#endif // __CUDACC__ +#endif // BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT #endif // BOOST_SAFE_NUMBERS_THROW_EXCEPTION_HPP diff --git a/include/boost/safe_numbers/detail/unsigned_integer_basis.hpp b/include/boost/safe_numbers/detail/unsigned_integer_basis.hpp index 10a97ae..cd8e1c1 100644 --- a/include/boost/safe_numbers/detail/unsigned_integer_basis.hpp +++ b/include/boost/safe_numbers/detail/unsigned_integer_basis.hpp @@ -427,7 +427,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr unsigned_integer_basis::oper namespace impl { -#if BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_add_overflow) && !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) +#if BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_add_overflow) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) template bool unsigned_intrin_add(const T lhs, const T rhs, T& result) @@ -435,7 +435,7 @@ bool unsigned_intrin_add(const T lhs, const T rhs, T& result) return __builtin_add_overflow(lhs, rhs, &result); } -#elif defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X64_INTRIN) && !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) +#elif defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X64_INTRIN) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) template bool unsigned_intrin_add(const T lhs, const T rhs, T& result) @@ -458,7 +458,7 @@ bool unsigned_intrin_add(const T lhs, const T rhs, T& result) } } -#elif defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN) && !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) +#elif defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) template bool unsigned_intrin_add(const T lhs, const T rhs, T& result) @@ -526,7 +526,7 @@ struct add_helper auto handle_overflow = [&res] { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -583,7 +583,7 @@ struct add_helper if constexpr (!std::is_same_v) { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (!std::is_constant_evaluated()) { @@ -644,7 +644,7 @@ struct add_helper if constexpr (!std::is_same_v) { - #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_add_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_addcarry_u64) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN)) && !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_add_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_addcarry_u64) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN)) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (!std::is_constant_evaluated()) { @@ -677,7 +677,7 @@ struct add_helper if constexpr (!std::is_same_v) { - #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_add_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_addcarry_u64) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN)) && !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_add_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_addcarry_u64) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN)) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (!std::is_constant_evaluated()) { @@ -763,7 +763,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE 1 error generated. */ - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { @@ -956,7 +956,7 @@ constexpr auto unsigned_integer_basis::operator+=(const unsigned_inte namespace impl { -#if BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_sub_overflow) && !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) +#if BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_sub_overflow) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) template bool unsigned_intrin_sub(T lhs, T rhs, T& result) @@ -964,7 +964,7 @@ bool unsigned_intrin_sub(T lhs, T rhs, T& result) return __builtin_sub_overflow(lhs, rhs, &result); } -#elif defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X64_INTRIN) && !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) +#elif defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X64_INTRIN) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) template bool unsigned_intrin_sub(T lhs, T rhs, T& result) @@ -987,7 +987,7 @@ bool unsigned_intrin_sub(T lhs, T rhs, T& result) } } -#elif defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN) && !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) +#elif defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) template bool unsigned_intrin_sub(T lhs, T rhs, T& result) @@ -1056,7 +1056,7 @@ struct sub_helper auto handle_underflow = [&res] { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { @@ -1108,7 +1108,7 @@ struct sub_helper if constexpr (!std::is_same_v) { - #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_sub_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_subborrow_u64) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN)) && !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_sub_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_subborrow_u64) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN)) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (!std::is_constant_evaluated()) { @@ -1169,7 +1169,7 @@ struct sub_helper if constexpr (!std::is_same_v) { - #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_sub_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_subborrow_u64) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN)) && !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_sub_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_subborrow_u64) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN)) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (!std::is_constant_evaluated()) { @@ -1202,7 +1202,7 @@ struct sub_helper if constexpr (!std::is_same_v) { - #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_sub_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_subborrow_u64) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN)) && !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_sub_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_subborrow_u64) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN)) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (!std::is_constant_evaluated()) { @@ -1232,7 +1232,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator-(const unsigned_integer_basis lhs, const unsigned_integer_basis rhs) -> unsigned_integer_basis { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { @@ -1287,7 +1287,7 @@ constexpr auto unsigned_integer_basis::operator-=(const unsigned_inte namespace impl { -#if BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_mul_overflow) && !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) +#if BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_mul_overflow) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) template bool unsigned_intrin_mul(const T lhs, const T rhs, T& result) @@ -1307,7 +1307,7 @@ inline bool unsigned_intrin_mul(const int128::uint128_t lhs, const int128::uint1 #endif -#elif defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X64_INTRIN) && !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) +#elif defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X64_INTRIN) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) template bool unsigned_intrin_mul(const T lhs, const T rhs, T& result) @@ -1331,7 +1331,7 @@ bool unsigned_intrin_mul(const T lhs, const T rhs, T& result) } } -#elif defined(BOOST_SAFENUMBERS_HAS_WINDOWS_ARM64_INTRIN) && !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) +#elif defined(BOOST_SAFENUMBERS_HAS_WINDOWS_ARM64_INTRIN) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) template bool unsigned_intrin_mul(const T lhs, const T rhs, T& result) @@ -1396,7 +1396,7 @@ struct mul_helper auto handle_overflow = [&res] { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { if constexpr (std::is_same_v) @@ -1450,7 +1450,7 @@ struct mul_helper if constexpr (!std::is_same_v) #endif { - #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_mul_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_umul128)) && !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_mul_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_umul128)) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (!std::is_constant_evaluated()) { @@ -1514,7 +1514,7 @@ struct mul_helper if constexpr (!std::is_same_v) #endif { - #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_mul_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_umul128)) && !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_mul_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_umul128)) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (!std::is_constant_evaluated()) { @@ -1550,7 +1550,7 @@ struct mul_helper if constexpr (!std::is_same_v) #endif { - #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_mul_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_umul128)) && !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_mul_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_umul128)) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (!std::is_constant_evaluated()) { @@ -1596,7 +1596,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator*(const unsigned_integer_basis lhs, const unsigned_integer_basis rhs) -> unsigned_integer_basis { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { @@ -1776,7 +1776,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator/(const unsigned_integer_basis lhs, const unsigned_integer_basis rhs) -> unsigned_integer_basis { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { @@ -1944,7 +1944,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator%(const unsigned_integer_basis lhs, const unsigned_integer_basis rhs) -> unsigned_integer_basis { - #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) + #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) if (std::is_constant_evaluated()) { diff --git a/include/boost/safe_numbers/device_error_reporting.hpp b/include/boost/safe_numbers/device_error_reporting.hpp new file mode 100644 index 0000000..e6e0011 --- /dev/null +++ b/include/boost/safe_numbers/device_error_reporting.hpp @@ -0,0 +1,28 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// Public entry point for device error reporting. Include this to obtain +// device_error_context. It pulls in the backend-neutral pieces (so +// device_exception_mode is always available) and dispatches to the CUDA or SYCL +// backend based on compiler detection and the enable macro. +// +// * CUDA: nvcc (__CUDACC__) + BOOST_SAFE_NUMBERS_ENABLE_CUDA +// * SYCL: BOOST_SAFE_NUMBERS_ENABLE_SYCL (with included first) +// +// On a pure host build neither backend is included and only device_exception_mode +// / the detail enums are defined. + +#ifndef BOOST_SAFE_NUMBERS_DEVICE_ERROR_REPORTING_HPP +#define BOOST_SAFE_NUMBERS_DEVICE_ERROR_REPORTING_HPP + +#include +#include + +#if defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA) +# include +#elif defined(BOOST_SAFE_NUMBERS_ENABLE_SYCL) +# include +#endif + +#endif // BOOST_SAFE_NUMBERS_DEVICE_ERROR_REPORTING_HPP diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 6a372f9..42ddea8 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -20,6 +20,27 @@ if(HAVE_BOOST_TEST) boost_test_jamfile(FILE cuda_jamfile LINK_LIBRARIES Boost::safe_numbers Boost::random CUDA::cudart COMPILE_DEFINITIONS BOOST_SAFE_NUMBERS_ENABLE_CUDA=1) + elseif(BOOST_SAFE_NUMBERS_ENABLE_SYCL) + + message(STATUS "Building Boost.safe_numbers with SYCL") + + set(CMAKE_CXX_COMPILER "icpx") + set(CMAKE_C_COMPILER "icx") + + # -fsycl must be on the link line too (it bundles the device image); COMPILE_OPTIONS + # only reaches the compile step, so add it to the linker flags as well. Without this + # every test links with no device image and segfaults at kernel submission. + set(CMAKE_EXE_LINKER_FLAGS "${CMAKE_EXE_LINKER_FLAGS} -fsycl") + + enable_testing() + + # __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__: the DPC++ range-rounding wrapper + # (used for a plain parallel_for over a non-round range) cannot materialize the + # small-integer llvm.umul.with.overflow intrinsic that unsigned multiplication of + # u8/u16/u32 lowers to, and aborts the kernel at JIT time. Disabling range rounding + # avoids the wrapper. See the note in doc/.../sycl.adoc. + boost_test_jamfile(FILE sycl_jamfile LINK_LIBRARIES Boost::safe_numbers Boost::charconv sycl COMPILE_DEFINITIONS BOOST_SAFE_NUMBERS_ENABLE_SYCL=1 __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__=1 COMPILE_OPTIONS -fsycl) + else() if(CMAKE_SYSTEM_PROCESSOR MATCHES "x86_64|AMD64|amd64|i386|i686|x86" AND NOT MINGW) diff --git a/test/sycl_jamfile b/test/sycl_jamfile new file mode 100644 index 0000000..c8838fd --- /dev/null +++ b/test/sycl_jamfile @@ -0,0 +1,114 @@ +# Copyright 2026 Matt Borland +# Distributed under the Boost Software License, Version 1.0. +# https://www.boost.org/LICENSE_1_0.txt +# +# SYCL test manifest. Consumed only by test/CMakeLists.txt via boost_test_jamfile +# under -DBOOST_SAFE_NUMBERS_ENABLE_SYCL=ON (icpx). Not built by plain b2. + +import testing ; +import ../../config/checks/config : requires ; + +project : requirements + [ requires cxx20_hdr_bit cxx20_hdr_compare cxx20_hdr_concepts ] + ; + +# Arithmetic - unsigned (u8, u16, u32, u64, u128) +run test_unsigned_add_sycl.cpp ; +run test_unsigned_sub_sycl.cpp ; +run test_unsigned_mul_sycl.cpp ; +run test_unsigned_div_sycl.cpp ; +run test_unsigned_mod_sycl.cpp ; + +# Arithmetic - signed (i8, i16, i32, i64, i128) +run test_signed_add_sycl.cpp ; +run test_signed_sub_sycl.cpp ; +run test_signed_mul_sycl.cpp ; +run test_signed_div_sycl.cpp ; +run test_signed_mod_sycl.cpp ; + +# Arithmetic - floats (f32, f64) +run test_float_add_sycl.cpp ; +run test_float_sub_sycl.cpp ; +run test_float_mul_sycl.cpp ; +run test_float_div_sycl.cpp ; + +# Bounded types (bounded_uint, bounded_int, bounded_float) +run test_bounded_add_sycl.cpp ; +run test_bounded_sub_sycl.cpp ; +run test_bounded_mul_sycl.cpp ; +run test_bounded_div_sycl.cpp ; +run test_bounded_ops_sycl.cpp ; + +# Error paths: the device records the error into the device_global buffer, +# and the host re-raises the matching std exception +run test_unsigned_add_error_sycl.cpp ; +run test_unsigned_sub_error_sycl.cpp ; +run test_unsigned_mul_error_sycl.cpp ; +run test_unsigned_div_error_sycl.cpp ; +run test_unsigned_mod_error_sycl.cpp ; +run test_signed_ops_error_sycl.cpp ; +run test_float_add_error_sycl.cpp ; +run test_float_sub_error_sycl.cpp ; +run test_float_mul_error_sycl.cpp ; +run test_float_div_error_sycl.cpp ; +run test_bounded_add_error_sycl.cpp ; + +# Error-reporting mechanism +run test_device_error_context_reuse_sycl.cpp ; +run test_device_error_single_context_sycl.cpp ; +run test_device_error_mapping_sycl.cpp ; +run test_device_error_workitem_id_sycl.cpp ; + +# Free functions - bit manipulation +run test_has_single_bit_sycl.cpp ; +run test_bit_ceil_sycl.cpp ; +run test_bit_floor_sycl.cpp ; +run test_bit_width_sycl.cpp ; +run test_bitswap_sycl.cpp ; +run test_byteswap_sycl.cpp ; +run test_countl_one_sycl.cpp ; +run test_countl_zero_sycl.cpp ; +run test_countr_one_sycl.cpp ; +run test_countr_zero_sycl.cpp ; +run test_popcount_sycl.cpp ; +run test_rotl_sycl.cpp ; +run test_rotr_sycl.cpp ; + +# Free functions - byte order conversions +run test_to_be_sycl.cpp ; +run test_from_be_sycl.cpp ; +run test_to_le_sycl.cpp ; +run test_from_le_sycl.cpp ; + +# Free functions - integer utilities +run test_isqrt_sycl.cpp ; +run test_ilog2_sycl.cpp ; +run test_ilog10_sycl.cpp ; +run test_ilog_sycl.cpp ; +run test_is_power_2_sycl.cpp ; +run test_is_power_10_sycl.cpp ; +run test_remove_trailing_zeros_sycl.cpp ; +run test_abs_diff_sycl.cpp ; +run test_div_ceil_sycl.cpp ; +run test_next_multiple_of_sycl.cpp ; +# test_ipow_sycl is disabled: its large looped kernels (u8..u128 in one TU) hit +# a DPC++ JIT bug where llvm.umul.with.overflow.i8/i16/i32 fail to materialize +# ("JIT session error: Symbols not found: [ old_llvm.umul.with.overflow.* ]", +# then abort). Reproduced identically on icpx 2026.0.0 (native OpenCL CPU +# runtime 2026.21.3) and 2026.1.0, with range rounding disabled; the plain +# unsigned mul tests that use the same intrinsic pass, so it is kernel-size +# dependent. Same class as the int128 signed-midpoint ICE; report to Intel. +# run test_ipow_sycl.cpp ; + +# Free functions - numeric algorithms +run test_gcd_sycl.cpp ; +run test_lcm_sycl.cpp ; +run test_midpoint_sycl.cpp ; + +# Free functions - charconv (to_chars / from_chars round-trip) +run test_charconv_sycl.cpp ; +run test_charconv_all_bases_sycl.cpp ; + +# Examples +run ../examples/sycl.cpp ; +run ../examples/sycl_error_handling.cpp ; diff --git a/test/sycl_test.hpp b/test/sycl_test.hpp new file mode 100644 index 0000000..a669bb7 --- /dev/null +++ b/test/sycl_test.hpp @@ -0,0 +1,399 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// Shared harness for the Boost.safe_numbers SYCL tests. Each test runs an +// operation element-wise on the SYCL device over random inputs and verifies +// that the device results match a host recomputation. The generic runners loop +// over every applicable width so one test file covers the whole type family. +// +// is included first so SYCL_EXTERNAL exists before the +// safe_numbers headers (required by BOOST_SAFE_NUMBERS_ENABLE_SYCL). + +#ifndef BOOST_SAFE_NUMBERS_TEST_SYCL_TEST_HPP +#define BOOST_SAFE_NUMBERS_TEST_SYCL_TEST_HPP + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include + +namespace sn_sycl_test { + +constexpr int num_elements {20000}; + +// Draw an operand whose magnitude occupies only the low half of the basis value +// bits (signed operands also get a random sign). Two such operands keep add, +// sub, mul, div and mod within range under the default throwing policy, so the +// arithmetic runners never trip the overflow machinery. +template +SafeT draw_half(std::mt19937_64& rng) +{ + using basis = typename SafeT::basis_type; + constexpr bool is_signed {std::numeric_limits::is_signed}; + constexpr int half_bits {is_signed ? (std::numeric_limits::digits / 2 - 1) + : (std::numeric_limits::digits / 2)}; + const std::uint64_t mask {half_bits >= 64 ? ~std::uint64_t{0} + : ((std::uint64_t{1} << half_bits) - 1)}; + const auto magnitude {static_cast(rng() & mask)}; + + if constexpr (is_signed) + { + return SafeT{(rng() & 1U) ? static_cast(-magnitude) : magnitude}; + } + else + { + return SafeT{magnitude}; + } +} + +// Draw a full-range operand (for bitwise ops, comparisons, byte swaps). +template +SafeT draw_full(std::mt19937_64& rng) +{ + using basis = typename SafeT::basis_type; + + if constexpr (sizeof(basis) > 8) + { + using high_type = decltype(basis{}.high); + return SafeT{basis{static_cast(rng()), static_cast(rng())}}; + } + else + { + return SafeT{static_cast(rng())}; + } +} + +// Divisor that avoids division by zero (magnitudes are already small enough to +// avoid the INT_MIN / -1 case). Callable on host and device. +template +SafeT nonzero(const SafeT b) noexcept +{ + return (b == SafeT{0}) ? SafeT{1} : b; +} + +// Uniform real operand in [lo, hi] for the float types (mirrors the ranges the +// CUDA float value tests use so every op stays finite). +template +SafeT draw_real(std::mt19937_64& rng, const double lo, const double hi) +{ + using basis = typename SafeT::basis_type; + std::uniform_real_distribution dist {static_cast(lo), static_cast(hi)}; + return SafeT{dist(rng)}; +} + +// Generic binary-op runner: op(a, b) is evaluated on the device for every +// element and re-evaluated on the host; any mismatch fails. op must depend only +// on device-enabled safe_numbers facilities. +template +int run_binary(const char* label, Gen gen, Op op) +{ + sycl::queue q; + + SafeT* a {sycl::malloc_shared(num_elements, q)}; + SafeT* b {sycl::malloc_shared(num_elements, q)}; + SafeT* out {sycl::malloc_shared(num_elements, q)}; + + std::mt19937_64 rng {42}; + for (int i {0}; i < num_elements; ++i) + { + a[i] = gen(rng); + b[i] = gen(rng); + } + + q.parallel_for(sycl::range<1>(num_elements), [=](sycl::id<1> idx) + { + const int i {static_cast(idx[0])}; + out[i] = op(a[i], b[i]); + }).wait(); + + int failures {0}; + for (int i {0}; i < num_elements; ++i) + { + if (!(out[i] == op(a[i], b[i]))) + { + ++failures; + } + } + + sycl::free(a, q); + sycl::free(b, q); + sycl::free(out, q); + + std::cout << label << ": " << (failures == 0 ? "PASSED" : "FAILED") << "\n"; + return failures == 0 ? EXIT_SUCCESS : EXIT_FAILURE; +} + +// Runs an arithmetic op (op-safe half-magnitude inputs) across every unsigned width. +template +int run_all_unsigned(const char* op_name, Op op) +{ + int rc {EXIT_SUCCESS}; + rc |= run_binary ("u8 " , &draw_half, op); + rc |= run_binary ("u16 " , &draw_half, op); + rc |= run_binary ("u32 " , &draw_half, op); + rc |= run_binary ("u64 " , &draw_half, op); + rc |= run_binary("u128 " , &draw_half, op); + static_cast(op_name); + return rc; +} + +// Runs an arithmetic op across every signed width. +template +int run_all_signed(const char* op_name, Op op) +{ + int rc {EXIT_SUCCESS}; + rc |= run_binary ("i8 " , &draw_half, op); + rc |= run_binary ("i16 " , &draw_half, op); + rc |= run_binary ("i32 " , &draw_half, op); + rc |= run_binary ("i64 " , &draw_half, op); + rc |= run_binary("i128 " , &draw_half, op); + static_cast(op_name); + return rc; +} + +// --------------------------------------------------------------------------- +// Free-function coverage (bit ops, integer utilities, numeric, byte conversions) +// --------------------------------------------------------------------------- + +// Input-generation strategy for a free-function runner, chosen so the tested +// operation stays well defined for every generated value under the default +// throwing policy (mirrors the input ranges the CUDA free-function tests use). +enum class gen_kind +{ + full, // [0, max] + nonzero, // [1, max] (ilog/ilog2/ilog10, is_power_10, remove_trailing_zeros) + below_half_max, // [0, max / 2] (bit_ceil, next_multiple_of first operand) + pos_10, // [1, 10] (lcm) + upto_10, // [0, 10] (ipow base) + upto_2, // [0, 2] (ipow exponent) +}; + +// Full-range value with the top value bit cleared, i.e. in [0, max / 2]. +template +SafeT draw_below_half_max(std::mt19937_64& rng) +{ + using basis = typename SafeT::basis_type; + return SafeT{static_cast(static_cast(draw_full(rng)) >> 1)}; +} + +// Small value in [lo, hi]. +template +SafeT draw_range(std::mt19937_64& rng, const std::uint64_t lo, const std::uint64_t hi) +{ + using basis = typename SafeT::basis_type; + return SafeT{static_cast(lo + (rng() % (hi - lo + 1)))}; +} + +template +SafeT gen_for(std::mt19937_64& rng) +{ + if constexpr (G == gen_kind::full) { return draw_full(rng); } + else if constexpr (G == gen_kind::nonzero) { return nonzero(draw_full(rng)); } + else if constexpr (G == gen_kind::below_half_max) { return draw_below_half_max(rng); } + else if constexpr (G == gen_kind::pos_10) { return draw_range(rng, 1, 10); } + else if constexpr (G == gen_kind::upto_10) { return draw_range(rng, 0, 10); } + else { return draw_range(rng, 0, 2); } +} + +// Generic unary free-function runner. The output type is deduced from the op, so +// it works whether the function returns the safe type, an int, or a bool. +template +int run_unary(const char* label, Gen gen, Op op) +{ + using OutT = std::decay_t()))>; + + sycl::queue q; + InT* in {sycl::malloc_shared(num_elements, q)}; + OutT* out {sycl::malloc_shared(num_elements, q)}; + + std::mt19937_64 rng {42}; + for (int i {0}; i < num_elements; ++i) + { + in[i] = gen(rng); + } + + q.parallel_for(sycl::range<1>(num_elements), [=](sycl::id<1> idx) + { + const int i {static_cast(idx[0])}; + out[i] = op(in[i]); + }).wait(); + + int failures {0}; + for (int i {0}; i < num_elements; ++i) + { + if (!(out[i] == op(in[i]))) + { + ++failures; + } + } + + sycl::free(in, q); + sycl::free(out, q); + + std::cout << label << ": " << (failures == 0 ? "PASSED" : "FAILED") << "\n"; + return failures == 0 ? EXIT_SUCCESS : EXIT_FAILURE; +} + +// Generic binary free-function runner with an independent generator per operand. +template +int run_binary_2(const char* label, GenA gen_a, GenB gen_b, Op op) +{ + using OutT = std::decay_t(), std::declval()))>; + + sycl::queue q; + InT* a {sycl::malloc_shared(num_elements, q)}; + InT* b {sycl::malloc_shared(num_elements, q)}; + OutT* out {sycl::malloc_shared(num_elements, q)}; + + std::mt19937_64 rng {42}; + for (int i {0}; i < num_elements; ++i) + { + a[i] = gen_a(rng); + b[i] = gen_b(rng); + } + + q.parallel_for(sycl::range<1>(num_elements), [=](sycl::id<1> idx) + { + const int i {static_cast(idx[0])}; + out[i] = op(a[i], b[i]); + }).wait(); + + int failures {0}; + for (int i {0}; i < num_elements; ++i) + { + if (!(out[i] == op(a[i], b[i]))) + { + ++failures; + } + } + + sycl::free(a, q); + sycl::free(b, q); + sycl::free(out, q); + + std::cout << label << ": " << (failures == 0 ? "PASSED" : "FAILED") << "\n"; + return failures == 0 ? EXIT_SUCCESS : EXIT_FAILURE; +} + +// Runs a unary free function across every unsigned width with the given generator. +template +int run_unary_all_unsigned(const char* name, Op op) +{ + int rc {EXIT_SUCCESS}; + rc |= run_unary (name, &gen_for, op); + rc |= run_unary (name, &gen_for, op); + rc |= run_unary (name, &gen_for, op); + rc |= run_unary (name, &gen_for, op); + rc |= run_unary(name, &gen_for, op); + return rc; +} + +// Runs a binary free function across every unsigned width, with an independent +// generator strategy for each operand. +template +int run_binary_all_unsigned(const char* name, Op op) +{ + int rc {EXIT_SUCCESS}; + rc |= run_binary_2 (name, &gen_for, &gen_for, op); + rc |= run_binary_2 (name, &gen_for, &gen_for, op); + rc |= run_binary_2 (name, &gen_for, &gen_for, op); + rc |= run_binary_2 (name, &gen_for, &gen_for, op); + rc |= run_binary_2(name, &gen_for, &gen_for, op); + return rc; +} + +// Launches `launch(q)` (which submits a kernel that reports an error for at +// least one element) inside a device_error_context and asserts synchronize() +// throws the expected std exception type. +template +int expect_device_throw(const char* label, Launch launch) +{ + sycl::queue q; + boost::safe_numbers::device_error_context ctx {q}; + + launch(q); + + try + { + ctx.synchronize(); + } + catch (const ExpectedException& e) + { + std::cout << label << ": PASSED (" << e.what() << ")\n"; + return EXIT_SUCCESS; + } + catch (const std::exception& e) + { + std::cerr << label << ": FAILED - wrong exception type: " << e.what() << "\n"; + return EXIT_FAILURE; + } + + std::cerr << label << ": FAILED - no exception thrown\n"; + return EXIT_FAILURE; +} + +// Fills two buffers with the given operands, applies the binary op on the +// device for every element (errors are recorded implicitly through the +// device_global buffer, exactly as on CUDA), and asserts synchronize() throws +// Expected. This is the one-liner behind the CUDA-parity error matrix. +template +int expect_binary_throw(const char* label, const SafeT a_val, const SafeT b_val, Op op) +{ + sycl::queue q; + boost::safe_numbers::device_error_context ctx {q}; + + constexpr int n {256}; + auto* a {sycl::malloc_shared(n, q)}; + auto* b {sycl::malloc_shared(n, q)}; + auto* out {sycl::malloc_shared(n, q)}; + for (int i {0}; i < n; ++i) + { + a[i] = a_val; + b[i] = b_val; + } + + q.parallel_for(sycl::range<1>(n), [=](sycl::id<1> idx) + { + const int i {static_cast(idx[0])}; + out[i] = op(a[i], b[i]); + }); + + int rc {EXIT_FAILURE}; + try + { + ctx.synchronize(); + std::cerr << label << ": FAILED - no exception thrown\n"; + } + catch (const ExpectedException& e) + { + std::cout << label << ": PASSED (" << e.what() << ")\n"; + rc = EXIT_SUCCESS; + } + catch (const std::exception& e) + { + std::cerr << label << ": FAILED - wrong exception type: " << e.what() << "\n"; + } + + sycl::free(a, q); + sycl::free(b, q); + sycl::free(out, q); + return rc; +} + +} // namespace sn_sycl_test + +#endif // BOOST_SAFE_NUMBERS_TEST_SYCL_TEST_HPP diff --git a/test/test_abs_diff_sycl.cpp b/test/test_abs_diff_sycl.cpp new file mode 100644 index 0000000..c97b042 --- /dev/null +++ b/test/test_abs_diff_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_binary_all_unsigned("abs_diff", [](auto a, auto b) { return boost::safe_numbers::abs_diff(a, b); }); +} diff --git a/test/test_bit_ceil_sycl.cpp b/test/test_bit_ceil_sycl.cpp new file mode 100644 index 0000000..7139586 --- /dev/null +++ b/test/test_bit_ceil_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("bit_ceil", [](auto a) { return boost::safe_numbers::bit_ceil(a); }); +} diff --git a/test/test_bit_floor_sycl.cpp b/test/test_bit_floor_sycl.cpp new file mode 100644 index 0000000..43e7d1d --- /dev/null +++ b/test/test_bit_floor_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("bit_floor", [](auto a) { return boost::safe_numbers::bit_floor(a); }); +} diff --git a/test/test_bit_width_sycl.cpp b/test/test_bit_width_sycl.cpp new file mode 100644 index 0000000..b9244ca --- /dev/null +++ b/test/test_bit_width_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("bit_width", [](auto a) { return boost::safe_numbers::bit_width(a); }); +} diff --git a/test/test_bitswap_sycl.cpp b/test/test_bitswap_sycl.cpp new file mode 100644 index 0000000..ec70f55 --- /dev/null +++ b/test/test_bitswap_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("bitswap", [](auto a) { return boost::safe_numbers::bitswap(a); }); +} diff --git a/test/test_bounded_add_error_sycl.cpp b/test/test_bounded_add_error_sycl.cpp new file mode 100644 index 0000000..51dd891 --- /dev/null +++ b/test/test_bounded_add_error_sycl.cpp @@ -0,0 +1,38 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// A bounded addition whose result leaves the declared range must be reported +// from the device and surface on the host as std::domain_error, for all +// three bounded families. + +#include "sycl_test.hpp" +#include +#include + +int main() +{ + using bu = boost::safe_numbers::bounded_uint<0U, 1000U>; + using bi = boost::safe_numbers::bounded_int<-1000, 1000>; + using bf = boost::safe_numbers::bounded_float<0.0F, 1000.0F>; + + int rc {EXIT_SUCCESS}; + + // 600 + 600 = 1200 exceeds every declared range below + rc |= sn_sycl_test::expect_binary_throw( + "bounded_uint add out of range", + bu{600U}, bu{600U}, + [](auto x, auto y) { return x + y; }); + + rc |= sn_sycl_test::expect_binary_throw( + "bounded_int add out of range", + bi{600}, bi{600}, + [](auto x, auto y) { return x + y; }); + + rc |= sn_sycl_test::expect_binary_throw( + "bounded_float add out of range", + bf{600.0F}, bf{600.0F}, + [](auto x, auto y) { return x + y; }); + + return rc; +} diff --git a/test/test_bounded_add_sycl.cpp b/test/test_bounded_add_sycl.cpp new file mode 100644 index 0000000..432eb1c --- /dev/null +++ b/test/test_bounded_add_sycl.cpp @@ -0,0 +1,47 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// In-range bounded addition on the device for all three bounded families, +// verified exactly against a host recomputation. Operand ranges keep every +// result inside the declared bounds. + +#include "sycl_test.hpp" +#include +#include + +using bu = boost::safe_numbers::bounded_uint<0U, 1000000U>; +using bi = boost::safe_numbers::bounded_int<-1000000, 1000000>; +using bf = boost::safe_numbers::bounded_float<0.0F, 1000000.0F>; + +// Uniform draw in [lo, hi] constructed through the bounded type's underlying type +template +BoundedT draw(std::mt19937_64& rng, const long long lo, const long long hi) +{ + const auto span {static_cast(hi - lo + 1)}; + return BoundedT{static_cast(lo + static_cast(rng() % span))}; +} + +int main() +{ + const auto op = [](auto a, auto b) { return a + b; }; + + int rc {EXIT_SUCCESS}; + + rc |= sn_sycl_test::run_binary_2("bounded_uint add", + [](std::mt19937_64& rng) { return draw(rng, 0, 400000); }, + [](std::mt19937_64& rng) { return draw(rng, 0, 400000); }, + op); + + rc |= sn_sycl_test::run_binary_2("bounded_int add", + [](std::mt19937_64& rng) { return draw(rng, -400000, 400000); }, + [](std::mt19937_64& rng) { return draw(rng, -400000, 400000); }, + op); + + rc |= sn_sycl_test::run_binary_2("bounded_float add", + [](std::mt19937_64& rng) { return draw(rng, 0, 400000); }, + [](std::mt19937_64& rng) { return draw(rng, 0, 400000); }, + op); + + return rc; +} diff --git a/test/test_bounded_div_sycl.cpp b/test/test_bounded_div_sycl.cpp new file mode 100644 index 0000000..994de35 --- /dev/null +++ b/test/test_bounded_div_sycl.cpp @@ -0,0 +1,47 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// In-range bounded division on the device for all three bounded families, +// verified exactly against a host recomputation. Operand ranges keep every +// result inside the declared bounds. + +#include "sycl_test.hpp" +#include +#include + +using bu = boost::safe_numbers::bounded_uint<0U, 1000000U>; +using bi = boost::safe_numbers::bounded_int<-1000000, 1000000>; +using bf = boost::safe_numbers::bounded_float<0.0F, 1000000.0F>; + +// Uniform draw in [lo, hi] constructed through the bounded type's underlying type +template +BoundedT draw(std::mt19937_64& rng, const long long lo, const long long hi) +{ + const auto span {static_cast(hi - lo + 1)}; + return BoundedT{static_cast(lo + static_cast(rng() % span))}; +} + +int main() +{ + const auto op = [](auto a, auto b) { return a / b; }; + + int rc {EXIT_SUCCESS}; + + rc |= sn_sycl_test::run_binary_2("bounded_uint div", + [](std::mt19937_64& rng) { return draw(rng, 0, 1000000); }, + [](std::mt19937_64& rng) { return draw(rng, 1, 1000); }, + op); + + rc |= sn_sycl_test::run_binary_2("bounded_int div", + [](std::mt19937_64& rng) { return draw(rng, -1000000, 1000000); }, + [](std::mt19937_64& rng) { return draw(rng, 1, 1000); }, + op); + + rc |= sn_sycl_test::run_binary_2("bounded_float div", + [](std::mt19937_64& rng) { return draw(rng, 0, 1000000); }, + [](std::mt19937_64& rng) { return draw(rng, 1, 1000); }, + op); + + return rc; +} diff --git a/test/test_bounded_mul_sycl.cpp b/test/test_bounded_mul_sycl.cpp new file mode 100644 index 0000000..bf121ff --- /dev/null +++ b/test/test_bounded_mul_sycl.cpp @@ -0,0 +1,47 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// In-range bounded multiplication on the device for all three bounded families, +// verified exactly against a host recomputation. Operand ranges keep every +// result inside the declared bounds. + +#include "sycl_test.hpp" +#include +#include + +using bu = boost::safe_numbers::bounded_uint<0U, 1000000U>; +using bi = boost::safe_numbers::bounded_int<-1000000, 1000000>; +using bf = boost::safe_numbers::bounded_float<0.0F, 1000000.0F>; + +// Uniform draw in [lo, hi] constructed through the bounded type's underlying type +template +BoundedT draw(std::mt19937_64& rng, const long long lo, const long long hi) +{ + const auto span {static_cast(hi - lo + 1)}; + return BoundedT{static_cast(lo + static_cast(rng() % span))}; +} + +int main() +{ + const auto op = [](auto a, auto b) { return a * b; }; + + int rc {EXIT_SUCCESS}; + + rc |= sn_sycl_test::run_binary_2("bounded_uint mul", + [](std::mt19937_64& rng) { return draw(rng, 0, 1000); }, + [](std::mt19937_64& rng) { return draw(rng, 0, 1000); }, + op); + + rc |= sn_sycl_test::run_binary_2("bounded_int mul", + [](std::mt19937_64& rng) { return draw(rng, -1000, 1000); }, + [](std::mt19937_64& rng) { return draw(rng, -1000, 1000); }, + op); + + rc |= sn_sycl_test::run_binary_2("bounded_float mul", + [](std::mt19937_64& rng) { return draw(rng, 0, 1000); }, + [](std::mt19937_64& rng) { return draw(rng, 0, 1000); }, + op); + + return rc; +} diff --git a/test/test_bounded_ops_sycl.cpp b/test/test_bounded_ops_sycl.cpp new file mode 100644 index 0000000..5be7fbe --- /dev/null +++ b/test/test_bounded_ops_sycl.cpp @@ -0,0 +1,65 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// Exercises the remaining bounded-type operators on the device (modulo, +// increment/decrement, unary plus/minus, and compound assignment) to +// confirm they are all callable from device code and produce correct +// results (the SYCL analog of the CUDA bounded_ops test). + +#include "sycl_test.hpp" +#include +#include + +using bu = boost::safe_numbers::bounded_uint<0U, 1000U>; +using bi = boost::safe_numbers::bounded_int<-1000, 1000>; +using bf = boost::safe_numbers::bounded_float<-1000.0F, 1000.0F>; + +int main() +{ + sycl::queue q; + + auto* ok {sycl::malloc_shared(1, q)}; + *ok = 0; + + q.single_task([=]() + { + // bounded_uint: modulo, increment, decrement, compound assignment + const bu a {100U}; + const bu b {30U}; + bu inc {a}; + ++inc; + inc--; + bu comp {a}; + comp += b; + comp -= b; + comp *= bu{2U}; + comp /= bu{2U}; + const bool r1 {(a % b == bu{10U}) && (inc == a) && (comp == a) && (a > b)}; + + // bounded_int: modulo, unary plus/minus, increment/decrement + const bi x {-100}; + const bi y {7}; + const bi neg {-y}; + bi xc {x}; + ++xc; + --xc; + const bool r2 {(x % y == bi{-2}) && (neg == bi{-7}) && (+x == x) && (xc == x)}; + + // bounded_float: compound assignment and comparison + const bf p {bf::basis_type{2.0F}}; + const bf v {bf::basis_type{3.0F}}; + bf pc {p}; + pc *= v; + pc /= v; + const bool r3 {(pc == p) && (p < v)}; + + *ok = (r1 && r2 && r3) ? 1 : 0; + }).wait(); + + const bool passed {*ok == 1}; + sycl::free(ok, q); + + std::cout << "bounded ops: " << (passed ? "PASSED" : "FAILED") << "\n"; + return passed ? EXIT_SUCCESS : EXIT_FAILURE; +} diff --git a/test/test_bounded_sub_sycl.cpp b/test/test_bounded_sub_sycl.cpp new file mode 100644 index 0000000..f81fff0 --- /dev/null +++ b/test/test_bounded_sub_sycl.cpp @@ -0,0 +1,47 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// In-range bounded subtraction on the device for all three bounded families, +// verified exactly against a host recomputation. Operand ranges keep every +// result inside the declared bounds. + +#include "sycl_test.hpp" +#include +#include + +using bu = boost::safe_numbers::bounded_uint<0U, 1000000U>; +using bi = boost::safe_numbers::bounded_int<-1000000, 1000000>; +using bf = boost::safe_numbers::bounded_float<0.0F, 1000000.0F>; + +// Uniform draw in [lo, hi] constructed through the bounded type's underlying type +template +BoundedT draw(std::mt19937_64& rng, const long long lo, const long long hi) +{ + const auto span {static_cast(hi - lo + 1)}; + return BoundedT{static_cast(lo + static_cast(rng() % span))}; +} + +int main() +{ + const auto op = [](auto a, auto b) { return a - b; }; + + int rc {EXIT_SUCCESS}; + + rc |= sn_sycl_test::run_binary_2("bounded_uint sub", + [](std::mt19937_64& rng) { return draw(rng, 500000, 1000000); }, + [](std::mt19937_64& rng) { return draw(rng, 0, 500000); }, + op); + + rc |= sn_sycl_test::run_binary_2("bounded_int sub", + [](std::mt19937_64& rng) { return draw(rng, -400000, 400000); }, + [](std::mt19937_64& rng) { return draw(rng, -400000, 400000); }, + op); + + rc |= sn_sycl_test::run_binary_2("bounded_float sub", + [](std::mt19937_64& rng) { return draw(rng, 500000, 1000000); }, + [](std::mt19937_64& rng) { return draw(rng, 0, 500000); }, + op); + + return rc; +} diff --git a/test/test_byteswap_sycl.cpp b/test/test_byteswap_sycl.cpp new file mode 100644 index 0000000..d414b86 --- /dev/null +++ b/test/test_byteswap_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("byteswap", [](auto a) { return boost::safe_numbers::byteswap(a); }); +} diff --git a/test/test_charconv_all_bases_sycl.cpp b/test/test_charconv_all_bases_sycl.cpp new file mode 100644 index 0000000..171d8f4 --- /dev/null +++ b/test/test_charconv_all_bases_sycl.cpp @@ -0,0 +1,96 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// to_chars/from_chars round-trip over every base 2..36 on the SYCL device +// across the unsigned widths (mirrors test_cuda__charconv_all_bases.cu). + +#include +#include +#include + +#include +#include +#include +#include + +namespace +{ + +constexpr int num_elements {20000}; + +template +T draw_full(std::mt19937_64& rng) +{ + using basis = typename T::basis_type; + if constexpr (sizeof(basis) > 8) + { + using high_type = decltype(basis{}.high); + return T{basis{static_cast(rng()), static_cast(rng())}}; + } + else + { + return T{static_cast(rng())}; + } +} + +template +int all_bases(const char* label) +{ + sycl::queue q; + T* in {sycl::malloc_shared(num_elements, q)}; + int* pass {sycl::malloc_shared(num_elements, q)}; + + std::mt19937_64 rng {42}; + for (int i {0}; i < num_elements; ++i) + { + in[i] = draw_full(rng); + } + + q.parallel_for(sycl::range<1>(num_elements), [=](sycl::id<1> idx) + { + const int i {static_cast(idx[0])}; + int pass_count {0}; + for (int base {2}; base <= 36; ++base) + { + char buf[256] {}; + const auto r {boost::charconv::to_chars(buf, buf + sizeof(buf), in[i], base)}; + T parsed {}; + boost::charconv::from_chars(buf, r.ptr, parsed, base); + if (parsed == in[i]) + { + ++pass_count; + } + } + pass[i] = pass_count; + }).wait(); + + int failures {0}; + for (int i {0}; i < num_elements; ++i) + { + if (pass[i] != 35) // bases 2..36 inclusive + { + ++failures; + } + } + + sycl::free(in, q); + sycl::free(pass, q); + + std::cout << label << ": " << (failures == 0 ? "PASSED" : "FAILED") << "\n"; + return failures == 0 ? EXIT_SUCCESS : EXIT_FAILURE; +} + +} // namespace + +int main() +{ + using namespace boost::safe_numbers; + int rc {EXIT_SUCCESS}; + rc |= all_bases ("u8"); + rc |= all_bases("u16"); + rc |= all_bases("u32"); + rc |= all_bases("u64"); + rc |= all_bases("u128"); + return rc; +} diff --git a/test/test_charconv_sycl.cpp b/test/test_charconv_sycl.cpp new file mode 100644 index 0000000..10cc781 --- /dev/null +++ b/test/test_charconv_sycl.cpp @@ -0,0 +1,89 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// Base-10 to_chars/from_chars round-trip on the SYCL device across the unsigned +// widths. charconv pulls in Boost.Charconv, so it is kept out of the shared +// harness and driven directly here (mirrors test_cuda__charconv.cu). + +#include +#include +#include + +#include +#include +#include +#include + +namespace +{ + +constexpr int num_elements {20000}; + +template +T draw_full(std::mt19937_64& rng) +{ + using basis = typename T::basis_type; + if constexpr (sizeof(basis) > 8) + { + using high_type = decltype(basis{}.high); + return T{basis{static_cast(rng()), static_cast(rng())}}; + } + else + { + return T{static_cast(rng())}; + } +} + +template +int roundtrip(const char* label) +{ + sycl::queue q; + T* in {sycl::malloc_shared(num_elements, q)}; + T* out {sycl::malloc_shared(num_elements, q)}; + + std::mt19937_64 rng {42}; + for (int i {0}; i < num_elements; ++i) + { + in[i] = draw_full(rng); + } + + q.parallel_for(sycl::range<1>(num_elements), [=](sycl::id<1> idx) + { + const int i {static_cast(idx[0])}; + char buf[64] {}; + const auto r {boost::charconv::to_chars(buf, buf + sizeof(buf), in[i])}; + T parsed {}; + boost::charconv::from_chars(buf, r.ptr, parsed); + out[i] = parsed; + }).wait(); + + int failures {0}; + for (int i {0}; i < num_elements; ++i) + { + if (!(out[i] == in[i])) + { + ++failures; + } + } + + sycl::free(in, q); + sycl::free(out, q); + + std::cout << label << ": " << (failures == 0 ? "PASSED" : "FAILED") << "\n"; + return failures == 0 ? EXIT_SUCCESS : EXIT_FAILURE; +} + +} // namespace + +int main() +{ + using namespace boost::safe_numbers; + int rc {EXIT_SUCCESS}; + rc |= roundtrip ("u8"); + rc |= roundtrip("u16"); + rc |= roundtrip("u32"); + rc |= roundtrip("u64"); + rc |= roundtrip("u128"); + return rc; +} diff --git a/test/test_countl_one_sycl.cpp b/test/test_countl_one_sycl.cpp new file mode 100644 index 0000000..d344a2f --- /dev/null +++ b/test/test_countl_one_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("countl_one", [](auto a) { return boost::safe_numbers::countl_one(a); }); +} diff --git a/test/test_countl_zero_sycl.cpp b/test/test_countl_zero_sycl.cpp new file mode 100644 index 0000000..1149a24 --- /dev/null +++ b/test/test_countl_zero_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("countl_zero", [](auto a) { return boost::safe_numbers::countl_zero(a); }); +} diff --git a/test/test_countr_one_sycl.cpp b/test/test_countr_one_sycl.cpp new file mode 100644 index 0000000..e9a99c8 --- /dev/null +++ b/test/test_countr_one_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("countr_one", [](auto a) { return boost::safe_numbers::countr_one(a); }); +} diff --git a/test/test_countr_zero_sycl.cpp b/test/test_countr_zero_sycl.cpp new file mode 100644 index 0000000..c89c93c --- /dev/null +++ b/test/test_countr_zero_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("countr_zero", [](auto a) { return boost::safe_numbers::countr_zero(a); }); +} diff --git a/test/test_cuda_bounded_float_add.cu b/test/test_cuda_bounded_float_add.cu index 12c8830..c851761 100644 --- a/test/test_cuda_bounded_float_add.cu +++ b/test/test_cuda_bounded_float_add.cu @@ -6,7 +6,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_bounded_float_add_error.cu b/test/test_cuda_bounded_float_add_error.cu index b87050d..5174a06 100644 --- a/test/test_cuda_bounded_float_add_error.cu +++ b/test/test_cuda_bounded_float_add_error.cu @@ -6,7 +6,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_bounded_float_div.cu b/test/test_cuda_bounded_float_div.cu index 16a9ad2..437a9a5 100644 --- a/test/test_cuda_bounded_float_div.cu +++ b/test/test_cuda_bounded_float_div.cu @@ -7,7 +7,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_bounded_float_mul.cu b/test/test_cuda_bounded_float_mul.cu index b054088..a6d8b73 100644 --- a/test/test_cuda_bounded_float_mul.cu +++ b/test/test_cuda_bounded_float_mul.cu @@ -6,7 +6,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_bounded_float_sub.cu b/test/test_cuda_bounded_float_sub.cu index 1a5fba1..80e5c65 100644 --- a/test/test_cuda_bounded_float_sub.cu +++ b/test/test_cuda_bounded_float_sub.cu @@ -7,7 +7,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_bounded_int_add.cu b/test/test_cuda_bounded_int_add.cu index 930a11b..f952f3c 100644 --- a/test/test_cuda_bounded_int_add.cu +++ b/test/test_cuda_bounded_int_add.cu @@ -7,7 +7,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_bounded_int_add_error.cu b/test/test_cuda_bounded_int_add_error.cu index db3a7a9..f9dc14e 100644 --- a/test/test_cuda_bounded_int_add_error.cu +++ b/test/test_cuda_bounded_int_add_error.cu @@ -6,7 +6,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_bounded_int_div.cu b/test/test_cuda_bounded_int_div.cu index a5cfa16..faf8e57 100644 --- a/test/test_cuda_bounded_int_div.cu +++ b/test/test_cuda_bounded_int_div.cu @@ -7,7 +7,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_bounded_int_mul.cu b/test/test_cuda_bounded_int_mul.cu index 19d058f..a981deb 100644 --- a/test/test_cuda_bounded_int_mul.cu +++ b/test/test_cuda_bounded_int_mul.cu @@ -7,7 +7,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_bounded_int_sub.cu b/test/test_cuda_bounded_int_sub.cu index 926389c..bb73ff1 100644 --- a/test/test_cuda_bounded_int_sub.cu +++ b/test/test_cuda_bounded_int_sub.cu @@ -7,7 +7,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_bounded_ops.cu b/test/test_cuda_bounded_ops.cu index ed5261e..c652f92 100644 --- a/test/test_cuda_bounded_ops.cu +++ b/test/test_cuda_bounded_ops.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_bounded_uint_add.cu b/test/test_cuda_bounded_uint_add.cu index 0c32f0c..b70aa35 100644 --- a/test/test_cuda_bounded_uint_add.cu +++ b/test/test_cuda_bounded_uint_add.cu @@ -7,7 +7,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_bounded_uint_add_error.cu b/test/test_cuda_bounded_uint_add_error.cu index 90108c6..f8e2210 100644 --- a/test/test_cuda_bounded_uint_add_error.cu +++ b/test/test_cuda_bounded_uint_add_error.cu @@ -6,7 +6,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_bounded_uint_div.cu b/test/test_cuda_bounded_uint_div.cu index 0f9fda2..acd8159 100644 --- a/test/test_cuda_bounded_uint_div.cu +++ b/test/test_cuda_bounded_uint_div.cu @@ -7,7 +7,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_bounded_uint_mul.cu b/test/test_cuda_bounded_uint_mul.cu index 1d3ae22..8a5656a 100644 --- a/test/test_cuda_bounded_uint_mul.cu +++ b/test/test_cuda_bounded_uint_mul.cu @@ -7,7 +7,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_bounded_uint_sub.cu b/test/test_cuda_bounded_uint_sub.cu index 61968fb..91dec8c 100644 --- a/test/test_cuda_bounded_uint_sub.cu +++ b/test/test_cuda_bounded_uint_sub.cu @@ -7,7 +7,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_error_handling.cu b/test/test_cuda_error_handling.cu index 5348609..ea70359 100644 --- a/test/test_cuda_error_handling.cu +++ b/test/test_cuda_error_handling.cu @@ -5,7 +5,7 @@ #include #include #include -#include +#include #include diff --git a/test/test_cuda_f32_add.cu b/test/test_cuda_f32_add.cu index c171564..26730cd 100644 --- a/test/test_cuda_f32_add.cu +++ b/test/test_cuda_f32_add.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_f32_add_error.cu b/test/test_cuda_f32_add_error.cu index d8d2ecf..92cee07 100644 --- a/test/test_cuda_f32_add_error.cu +++ b/test/test_cuda_f32_add_error.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_f32_div.cu b/test/test_cuda_f32_div.cu index ef5d8da..91a3137 100644 --- a/test/test_cuda_f32_div.cu +++ b/test/test_cuda_f32_div.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_f32_div_error.cu b/test/test_cuda_f32_div_error.cu index b761742..f59c2b8 100644 --- a/test/test_cuda_f32_div_error.cu +++ b/test/test_cuda_f32_div_error.cu @@ -8,7 +8,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_f32_mul.cu b/test/test_cuda_f32_mul.cu index 99a4ef5..7ba98b1 100644 --- a/test/test_cuda_f32_mul.cu +++ b/test/test_cuda_f32_mul.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_f32_mul_error.cu b/test/test_cuda_f32_mul_error.cu index 0db00bd..1c2a83f 100644 --- a/test/test_cuda_f32_mul_error.cu +++ b/test/test_cuda_f32_mul_error.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_f32_sub.cu b/test/test_cuda_f32_sub.cu index 75d060e..01eb4ec 100644 --- a/test/test_cuda_f32_sub.cu +++ b/test/test_cuda_f32_sub.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_f32_sub_error.cu b/test/test_cuda_f32_sub_error.cu index 82293b2..92c8dee 100644 --- a/test/test_cuda_f32_sub_error.cu +++ b/test/test_cuda_f32_sub_error.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_f64_add.cu b/test/test_cuda_f64_add.cu index 516f2a1..159b8d8 100644 --- a/test/test_cuda_f64_add.cu +++ b/test/test_cuda_f64_add.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_f64_add_error.cu b/test/test_cuda_f64_add_error.cu index 197e598..b72b6a0 100644 --- a/test/test_cuda_f64_add_error.cu +++ b/test/test_cuda_f64_add_error.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_f64_div.cu b/test/test_cuda_f64_div.cu index fafa358..b5c5419 100644 --- a/test/test_cuda_f64_div.cu +++ b/test/test_cuda_f64_div.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_f64_div_error.cu b/test/test_cuda_f64_div_error.cu index cebe472..cfe33f9 100644 --- a/test/test_cuda_f64_div_error.cu +++ b/test/test_cuda_f64_div_error.cu @@ -8,7 +8,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_f64_mul.cu b/test/test_cuda_f64_mul.cu index 2e3fc2a..29fca0b 100644 --- a/test/test_cuda_f64_mul.cu +++ b/test/test_cuda_f64_mul.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_f64_mul_error.cu b/test/test_cuda_f64_mul_error.cu index 13f950b..7687fcc 100644 --- a/test/test_cuda_f64_mul_error.cu +++ b/test/test_cuda_f64_mul_error.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_f64_sub.cu b/test/test_cuda_f64_sub.cu index 2f1430f..f4d6b6d 100644 --- a/test/test_cuda_f64_sub.cu +++ b/test/test_cuda_f64_sub.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_f64_sub_error.cu b/test/test_cuda_f64_sub_error.cu index a40dd7f..d3e1fc7 100644 --- a/test/test_cuda_f64_sub_error.cu +++ b/test/test_cuda_f64_sub_error.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_i128_add.cu b/test/test_cuda_i128_add.cu index ca3393a..1540991 100644 --- a/test/test_cuda_i128_add.cu +++ b/test/test_cuda_i128_add.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_i128_div.cu b/test/test_cuda_i128_div.cu index 885dfe2..601ca01 100644 --- a/test/test_cuda_i128_div.cu +++ b/test/test_cuda_i128_div.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_i128_mod.cu b/test/test_cuda_i128_mod.cu index 86248f3..a7e7e5b 100644 --- a/test/test_cuda_i128_mod.cu +++ b/test/test_cuda_i128_mod.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_i128_mul.cu b/test/test_cuda_i128_mul.cu index d64b1d5..57fa92c 100644 --- a/test/test_cuda_i128_mul.cu +++ b/test/test_cuda_i128_mul.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_i128_sub.cu b/test/test_cuda_i128_sub.cu index c6c4a61..fb91cf8 100644 --- a/test/test_cuda_i128_sub.cu +++ b/test/test_cuda_i128_sub.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_i16_add.cu b/test/test_cuda_i16_add.cu index d2869e8..6a666e8 100644 --- a/test/test_cuda_i16_add.cu +++ b/test/test_cuda_i16_add.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_i16_div.cu b/test/test_cuda_i16_div.cu index 1521140..746f587 100644 --- a/test/test_cuda_i16_div.cu +++ b/test/test_cuda_i16_div.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_i16_mod.cu b/test/test_cuda_i16_mod.cu index 2493d8a..46a7f73 100644 --- a/test/test_cuda_i16_mod.cu +++ b/test/test_cuda_i16_mod.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_i16_mul.cu b/test/test_cuda_i16_mul.cu index 0265be1..ec90117 100644 --- a/test/test_cuda_i16_mul.cu +++ b/test/test_cuda_i16_mul.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_i16_sub.cu b/test/test_cuda_i16_sub.cu index d867aed..e7d42e9 100644 --- a/test/test_cuda_i16_sub.cu +++ b/test/test_cuda_i16_sub.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_i32_add.cu b/test/test_cuda_i32_add.cu index ed0745d..b67efaf 100644 --- a/test/test_cuda_i32_add.cu +++ b/test/test_cuda_i32_add.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_i32_div.cu b/test/test_cuda_i32_div.cu index f18a61f..1f1ea71 100644 --- a/test/test_cuda_i32_div.cu +++ b/test/test_cuda_i32_div.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_i32_mod.cu b/test/test_cuda_i32_mod.cu index 7269b36..b816787 100644 --- a/test/test_cuda_i32_mod.cu +++ b/test/test_cuda_i32_mod.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_i32_mul.cu b/test/test_cuda_i32_mul.cu index 63d1465..444114d 100644 --- a/test/test_cuda_i32_mul.cu +++ b/test/test_cuda_i32_mul.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_i32_sub.cu b/test/test_cuda_i32_sub.cu index f67c522..4da0b94 100644 --- a/test/test_cuda_i32_sub.cu +++ b/test/test_cuda_i32_sub.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_i64_add.cu b/test/test_cuda_i64_add.cu index 98ef840..c9910d3 100644 --- a/test/test_cuda_i64_add.cu +++ b/test/test_cuda_i64_add.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_i64_div.cu b/test/test_cuda_i64_div.cu index 2a93597..6ad040b 100644 --- a/test/test_cuda_i64_div.cu +++ b/test/test_cuda_i64_div.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_i64_mod.cu b/test/test_cuda_i64_mod.cu index a17f228..041dad7 100644 --- a/test/test_cuda_i64_mod.cu +++ b/test/test_cuda_i64_mod.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_i64_mul.cu b/test/test_cuda_i64_mul.cu index 361f159..c9d1695 100644 --- a/test/test_cuda_i64_mul.cu +++ b/test/test_cuda_i64_mul.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_i64_sub.cu b/test/test_cuda_i64_sub.cu index 210624a..113b199 100644 --- a/test/test_cuda_i64_sub.cu +++ b/test/test_cuda_i64_sub.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_i8_add.cu b/test/test_cuda_i8_add.cu index 71fe64b..07b4a4a 100644 --- a/test/test_cuda_i8_add.cu +++ b/test/test_cuda_i8_add.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_i8_div.cu b/test/test_cuda_i8_div.cu index 0351aff..97ab720 100644 --- a/test/test_cuda_i8_div.cu +++ b/test/test_cuda_i8_div.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_i8_mod.cu b/test/test_cuda_i8_mod.cu index cb3b118..ead37d1 100644 --- a/test/test_cuda_i8_mod.cu +++ b/test/test_cuda_i8_mod.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_i8_mul.cu b/test/test_cuda_i8_mul.cu index f274dc6..49eb055 100644 --- a/test/test_cuda_i8_mul.cu +++ b/test/test_cuda_i8_mul.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_i8_sub.cu b/test/test_cuda_i8_sub.cu index 02bf40c..a5429b1 100644 --- a/test/test_cuda_i8_sub.cu +++ b/test/test_cuda_i8_sub.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_invalid_argument.cu b/test/test_cuda_invalid_argument.cu index 504937d..d65d0f7 100644 --- a/test/test_cuda_invalid_argument.cu +++ b/test/test_cuda_invalid_argument.cu @@ -12,7 +12,7 @@ #include #include #include -#include +#include #include diff --git a/test/test_cuda_u128_abs_diff.cu b/test/test_cuda_u128_abs_diff.cu index 1fa37dd..8e2312a 100644 --- a/test/test_cuda_u128_abs_diff.cu +++ b/test/test_cuda_u128_abs_diff.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_add.cu b/test/test_cuda_u128_add.cu index 251d15e..048930d 100644 --- a/test/test_cuda_u128_add.cu +++ b/test/test_cuda_u128_add.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_add_error.cu b/test/test_cuda_u128_add_error.cu index 4857ea9..a485c14 100644 --- a/test/test_cuda_u128_add_error.cu +++ b/test/test_cuda_u128_add_error.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u128_bit_ceil.cu b/test/test_cuda_u128_bit_ceil.cu index 5cfbcd2..160234f 100644 --- a/test/test_cuda_u128_bit_ceil.cu +++ b/test/test_cuda_u128_bit_ceil.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_bit_floor.cu b/test/test_cuda_u128_bit_floor.cu index 2cd917b..e8c96f5 100644 --- a/test/test_cuda_u128_bit_floor.cu +++ b/test/test_cuda_u128_bit_floor.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_bit_width.cu b/test/test_cuda_u128_bit_width.cu index 93a1574..7247b14 100644 --- a/test/test_cuda_u128_bit_width.cu +++ b/test/test_cuda_u128_bit_width.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_bitswap.cu b/test/test_cuda_u128_bitswap.cu index fee52ba..039f630 100644 --- a/test/test_cuda_u128_bitswap.cu +++ b/test/test_cuda_u128_bitswap.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_byteswap.cu b/test/test_cuda_u128_byteswap.cu index 8722bbf..9e6969c 100644 --- a/test/test_cuda_u128_byteswap.cu +++ b/test/test_cuda_u128_byteswap.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_charconv.cu b/test/test_cuda_u128_charconv.cu index 4dc7a5e..e9a5618 100644 --- a/test/test_cuda_u128_charconv.cu +++ b/test/test_cuda_u128_charconv.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_charconv_all_bases.cu b/test/test_cuda_u128_charconv_all_bases.cu index d5f06a5..2fcc476 100644 --- a/test/test_cuda_u128_charconv_all_bases.cu +++ b/test/test_cuda_u128_charconv_all_bases.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_countl_one.cu b/test/test_cuda_u128_countl_one.cu index a5ba398..cf6cdb2 100644 --- a/test/test_cuda_u128_countl_one.cu +++ b/test/test_cuda_u128_countl_one.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_countl_zero.cu b/test/test_cuda_u128_countl_zero.cu index f579b8a..885cee9 100644 --- a/test/test_cuda_u128_countl_zero.cu +++ b/test/test_cuda_u128_countl_zero.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_countr_one.cu b/test/test_cuda_u128_countr_one.cu index 2f937cf..73bc2b9 100644 --- a/test/test_cuda_u128_countr_one.cu +++ b/test/test_cuda_u128_countr_one.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_countr_zero.cu b/test/test_cuda_u128_countr_zero.cu index 3a703e7..13d58e0 100644 --- a/test/test_cuda_u128_countr_zero.cu +++ b/test/test_cuda_u128_countr_zero.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_div.cu b/test/test_cuda_u128_div.cu index 58959c4..93f9ace 100644 --- a/test/test_cuda_u128_div.cu +++ b/test/test_cuda_u128_div.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_div_ceil.cu b/test/test_cuda_u128_div_ceil.cu index 7687866..46af7ff 100644 --- a/test/test_cuda_u128_div_ceil.cu +++ b/test/test_cuda_u128_div_ceil.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_div_error.cu b/test/test_cuda_u128_div_error.cu index fa43fb3..f20b0ad 100644 --- a/test/test_cuda_u128_div_error.cu +++ b/test/test_cuda_u128_div_error.cu @@ -8,7 +8,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u128_from_be.cu b/test/test_cuda_u128_from_be.cu index a906677..e7686ff 100644 --- a/test/test_cuda_u128_from_be.cu +++ b/test/test_cuda_u128_from_be.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_from_le.cu b/test/test_cuda_u128_from_le.cu index 54e61e2..26a4d26 100644 --- a/test/test_cuda_u128_from_le.cu +++ b/test/test_cuda_u128_from_le.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_gcd.cu b/test/test_cuda_u128_gcd.cu index 45123ac..a7064ad 100644 --- a/test/test_cuda_u128_gcd.cu +++ b/test/test_cuda_u128_gcd.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_has_single_bit.cu b/test/test_cuda_u128_has_single_bit.cu index 3f4c2ba..9372541 100644 --- a/test/test_cuda_u128_has_single_bit.cu +++ b/test/test_cuda_u128_has_single_bit.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_ilog.cu b/test/test_cuda_u128_ilog.cu index 88b6741..8f3a131 100644 --- a/test/test_cuda_u128_ilog.cu +++ b/test/test_cuda_u128_ilog.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_ilog10.cu b/test/test_cuda_u128_ilog10.cu index 2296c6e..dacee20 100644 --- a/test/test_cuda_u128_ilog10.cu +++ b/test/test_cuda_u128_ilog10.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_ilog2.cu b/test/test_cuda_u128_ilog2.cu index 873423b..3429c74 100644 --- a/test/test_cuda_u128_ilog2.cu +++ b/test/test_cuda_u128_ilog2.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_ipow.cu b/test/test_cuda_u128_ipow.cu index 3e4225a..d9467c3 100644 --- a/test/test_cuda_u128_ipow.cu +++ b/test/test_cuda_u128_ipow.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_is_power_10.cu b/test/test_cuda_u128_is_power_10.cu index 94e8ec1..c37a27d 100644 --- a/test/test_cuda_u128_is_power_10.cu +++ b/test/test_cuda_u128_is_power_10.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_is_power_2.cu b/test/test_cuda_u128_is_power_2.cu index 6337c96..22b559d 100644 --- a/test/test_cuda_u128_is_power_2.cu +++ b/test/test_cuda_u128_is_power_2.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_isqrt.cu b/test/test_cuda_u128_isqrt.cu index 6c9ed6f..a35490b 100644 --- a/test/test_cuda_u128_isqrt.cu +++ b/test/test_cuda_u128_isqrt.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_lcm.cu b/test/test_cuda_u128_lcm.cu index 9e10fd6..3f653b9 100644 --- a/test/test_cuda_u128_lcm.cu +++ b/test/test_cuda_u128_lcm.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_midpoint.cu b/test/test_cuda_u128_midpoint.cu index 7c33fd9..cb14995 100644 --- a/test/test_cuda_u128_midpoint.cu +++ b/test/test_cuda_u128_midpoint.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_mod.cu b/test/test_cuda_u128_mod.cu index cbb2a6f..892136d 100644 --- a/test/test_cuda_u128_mod.cu +++ b/test/test_cuda_u128_mod.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_mod_error.cu b/test/test_cuda_u128_mod_error.cu index 3617b09..d89690a 100644 --- a/test/test_cuda_u128_mod_error.cu +++ b/test/test_cuda_u128_mod_error.cu @@ -6,7 +6,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u128_mul.cu b/test/test_cuda_u128_mul.cu index c74ea05..cf10f64 100644 --- a/test/test_cuda_u128_mul.cu +++ b/test/test_cuda_u128_mul.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_mul_error.cu b/test/test_cuda_u128_mul_error.cu index be483eb..a9d485c 100644 --- a/test/test_cuda_u128_mul_error.cu +++ b/test/test_cuda_u128_mul_error.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u128_next_multiple_of.cu b/test/test_cuda_u128_next_multiple_of.cu index 790bde6..1eca235 100644 --- a/test/test_cuda_u128_next_multiple_of.cu +++ b/test/test_cuda_u128_next_multiple_of.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_popcount.cu b/test/test_cuda_u128_popcount.cu index 8967452..078e2c3 100644 --- a/test/test_cuda_u128_popcount.cu +++ b/test/test_cuda_u128_popcount.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_remove_trailing_zeros.cu b/test/test_cuda_u128_remove_trailing_zeros.cu index fa803f5..c186c51 100644 --- a/test/test_cuda_u128_remove_trailing_zeros.cu +++ b/test/test_cuda_u128_remove_trailing_zeros.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_rotl.cu b/test/test_cuda_u128_rotl.cu index 37447f0..0869688 100644 --- a/test/test_cuda_u128_rotl.cu +++ b/test/test_cuda_u128_rotl.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_rotr.cu b/test/test_cuda_u128_rotr.cu index 6683a7a..a309b18 100644 --- a/test/test_cuda_u128_rotr.cu +++ b/test/test_cuda_u128_rotr.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_sub.cu b/test/test_cuda_u128_sub.cu index a2f50a9..0aaa00d 100644 --- a/test/test_cuda_u128_sub.cu +++ b/test/test_cuda_u128_sub.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_sub_error.cu b/test/test_cuda_u128_sub_error.cu index 2cdada8..b3c38f1 100644 --- a/test/test_cuda_u128_sub_error.cu +++ b/test/test_cuda_u128_sub_error.cu @@ -8,7 +8,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u128_to_be.cu b/test/test_cuda_u128_to_be.cu index 8fb1328..f0b58e1 100644 --- a/test/test_cuda_u128_to_be.cu +++ b/test/test_cuda_u128_to_be.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u128_to_le.cu b/test/test_cuda_u128_to_le.cu index 9aa25ae..bbda3b2 100644 --- a/test/test_cuda_u128_to_le.cu +++ b/test/test_cuda_u128_to_le.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include #include #include "cuda_managed_ptr.hpp" diff --git a/test/test_cuda_u16_abs_diff.cu b/test/test_cuda_u16_abs_diff.cu index a106567..460256a 100644 --- a/test/test_cuda_u16_abs_diff.cu +++ b/test/test_cuda_u16_abs_diff.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_add.cu b/test/test_cuda_u16_add.cu index 4ef66c9..f6e9479 100644 --- a/test/test_cuda_u16_add.cu +++ b/test/test_cuda_u16_add.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_add_error.cu b/test/test_cuda_u16_add_error.cu index b9e018a..d0409db 100644 --- a/test/test_cuda_u16_add_error.cu +++ b/test/test_cuda_u16_add_error.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u16_bit_ceil.cu b/test/test_cuda_u16_bit_ceil.cu index 063446b..9024a15 100644 --- a/test/test_cuda_u16_bit_ceil.cu +++ b/test/test_cuda_u16_bit_ceil.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_bit_floor.cu b/test/test_cuda_u16_bit_floor.cu index 7c28294..1d29e2f 100644 --- a/test/test_cuda_u16_bit_floor.cu +++ b/test/test_cuda_u16_bit_floor.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_bit_width.cu b/test/test_cuda_u16_bit_width.cu index e5c2634..2112a9a 100644 --- a/test/test_cuda_u16_bit_width.cu +++ b/test/test_cuda_u16_bit_width.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_bitswap.cu b/test/test_cuda_u16_bitswap.cu index f61d4b8..737af01 100644 --- a/test/test_cuda_u16_bitswap.cu +++ b/test/test_cuda_u16_bitswap.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_byteswap.cu b/test/test_cuda_u16_byteswap.cu index 6b0a89c..b6dfcb3 100644 --- a/test/test_cuda_u16_byteswap.cu +++ b/test/test_cuda_u16_byteswap.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_charconv.cu b/test/test_cuda_u16_charconv.cu index b503c6d..2c9bb76 100644 --- a/test/test_cuda_u16_charconv.cu +++ b/test/test_cuda_u16_charconv.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_charconv_all_bases.cu b/test/test_cuda_u16_charconv_all_bases.cu index 72c5d9f..8282970 100644 --- a/test/test_cuda_u16_charconv_all_bases.cu +++ b/test/test_cuda_u16_charconv_all_bases.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_countl_one.cu b/test/test_cuda_u16_countl_one.cu index f951719..9085ab6 100644 --- a/test/test_cuda_u16_countl_one.cu +++ b/test/test_cuda_u16_countl_one.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_countl_zero.cu b/test/test_cuda_u16_countl_zero.cu index 4879a07..7994932 100644 --- a/test/test_cuda_u16_countl_zero.cu +++ b/test/test_cuda_u16_countl_zero.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_countr_one.cu b/test/test_cuda_u16_countr_one.cu index 4bfce26..b0f7563 100644 --- a/test/test_cuda_u16_countr_one.cu +++ b/test/test_cuda_u16_countr_one.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_countr_zero.cu b/test/test_cuda_u16_countr_zero.cu index 87d2089..eed48f3 100644 --- a/test/test_cuda_u16_countr_zero.cu +++ b/test/test_cuda_u16_countr_zero.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_div.cu b/test/test_cuda_u16_div.cu index ec677da..574f41e 100644 --- a/test/test_cuda_u16_div.cu +++ b/test/test_cuda_u16_div.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_div_ceil.cu b/test/test_cuda_u16_div_ceil.cu index f0abefc..9fc1e3c 100644 --- a/test/test_cuda_u16_div_ceil.cu +++ b/test/test_cuda_u16_div_ceil.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_div_error.cu b/test/test_cuda_u16_div_error.cu index f3bf7a1..8ff8117 100644 --- a/test/test_cuda_u16_div_error.cu +++ b/test/test_cuda_u16_div_error.cu @@ -8,7 +8,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u16_from_be.cu b/test/test_cuda_u16_from_be.cu index 11ed7e3..0e4b1b8 100644 --- a/test/test_cuda_u16_from_be.cu +++ b/test/test_cuda_u16_from_be.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_from_le.cu b/test/test_cuda_u16_from_le.cu index 9058d38..3945b9d 100644 --- a/test/test_cuda_u16_from_le.cu +++ b/test/test_cuda_u16_from_le.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_gcd.cu b/test/test_cuda_u16_gcd.cu index 7ee5987..e1c28fb 100644 --- a/test/test_cuda_u16_gcd.cu +++ b/test/test_cuda_u16_gcd.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_has_single_bit.cu b/test/test_cuda_u16_has_single_bit.cu index 83a7929..ebf2e82 100644 --- a/test/test_cuda_u16_has_single_bit.cu +++ b/test/test_cuda_u16_has_single_bit.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_ilog.cu b/test/test_cuda_u16_ilog.cu index 94e336a..6f433ab 100644 --- a/test/test_cuda_u16_ilog.cu +++ b/test/test_cuda_u16_ilog.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_ilog10.cu b/test/test_cuda_u16_ilog10.cu index 2421729..35a8b67 100644 --- a/test/test_cuda_u16_ilog10.cu +++ b/test/test_cuda_u16_ilog10.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_ilog2.cu b/test/test_cuda_u16_ilog2.cu index 73ad788..8fd75ea 100644 --- a/test/test_cuda_u16_ilog2.cu +++ b/test/test_cuda_u16_ilog2.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_ipow.cu b/test/test_cuda_u16_ipow.cu index c5eea4e..58397ee 100644 --- a/test/test_cuda_u16_ipow.cu +++ b/test/test_cuda_u16_ipow.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_is_power_10.cu b/test/test_cuda_u16_is_power_10.cu index e9bf609..2409326 100644 --- a/test/test_cuda_u16_is_power_10.cu +++ b/test/test_cuda_u16_is_power_10.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_is_power_2.cu b/test/test_cuda_u16_is_power_2.cu index ccc59d9..db11139 100644 --- a/test/test_cuda_u16_is_power_2.cu +++ b/test/test_cuda_u16_is_power_2.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_isqrt.cu b/test/test_cuda_u16_isqrt.cu index 215374f..aad0e9c 100644 --- a/test/test_cuda_u16_isqrt.cu +++ b/test/test_cuda_u16_isqrt.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_lcm.cu b/test/test_cuda_u16_lcm.cu index 2862c58..6635e9f 100644 --- a/test/test_cuda_u16_lcm.cu +++ b/test/test_cuda_u16_lcm.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_midpoint.cu b/test/test_cuda_u16_midpoint.cu index acca756..6af9ee9 100644 --- a/test/test_cuda_u16_midpoint.cu +++ b/test/test_cuda_u16_midpoint.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_mod.cu b/test/test_cuda_u16_mod.cu index 489d47b..4900af2 100644 --- a/test/test_cuda_u16_mod.cu +++ b/test/test_cuda_u16_mod.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_mod_error.cu b/test/test_cuda_u16_mod_error.cu index 8361dec..179cf90 100644 --- a/test/test_cuda_u16_mod_error.cu +++ b/test/test_cuda_u16_mod_error.cu @@ -8,7 +8,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u16_mul.cu b/test/test_cuda_u16_mul.cu index ed83522..b24f0aa 100644 --- a/test/test_cuda_u16_mul.cu +++ b/test/test_cuda_u16_mul.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_mul_error.cu b/test/test_cuda_u16_mul_error.cu index 735ca7c..c9d1fef 100644 --- a/test/test_cuda_u16_mul_error.cu +++ b/test/test_cuda_u16_mul_error.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u16_next_multiple_of.cu b/test/test_cuda_u16_next_multiple_of.cu index 7a0d182..0d914bc 100644 --- a/test/test_cuda_u16_next_multiple_of.cu +++ b/test/test_cuda_u16_next_multiple_of.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_popcount.cu b/test/test_cuda_u16_popcount.cu index a55aa9d..c529011 100644 --- a/test/test_cuda_u16_popcount.cu +++ b/test/test_cuda_u16_popcount.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_remove_trailing_zeros.cu b/test/test_cuda_u16_remove_trailing_zeros.cu index db9afee..dca46ec 100644 --- a/test/test_cuda_u16_remove_trailing_zeros.cu +++ b/test/test_cuda_u16_remove_trailing_zeros.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_rotl.cu b/test/test_cuda_u16_rotl.cu index 3196b99..938107c 100644 --- a/test/test_cuda_u16_rotl.cu +++ b/test/test_cuda_u16_rotl.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_rotr.cu b/test/test_cuda_u16_rotr.cu index 94008f5..b5a2b67 100644 --- a/test/test_cuda_u16_rotr.cu +++ b/test/test_cuda_u16_rotr.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_sub.cu b/test/test_cuda_u16_sub.cu index d3aeffa..3e6993a 100644 --- a/test/test_cuda_u16_sub.cu +++ b/test/test_cuda_u16_sub.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_sub_error.cu b/test/test_cuda_u16_sub_error.cu index e86155c..301d7d2 100644 --- a/test/test_cuda_u16_sub_error.cu +++ b/test/test_cuda_u16_sub_error.cu @@ -8,7 +8,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u16_to_be.cu b/test/test_cuda_u16_to_be.cu index e7c7b3a..59265b1 100644 --- a/test/test_cuda_u16_to_be.cu +++ b/test/test_cuda_u16_to_be.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u16_to_le.cu b/test/test_cuda_u16_to_le.cu index 7f4e004..407ff6a 100644 --- a/test/test_cuda_u16_to_le.cu +++ b/test/test_cuda_u16_to_le.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_abs_diff.cu b/test/test_cuda_u32_abs_diff.cu index 2ffa845..f559775 100644 --- a/test/test_cuda_u32_abs_diff.cu +++ b/test/test_cuda_u32_abs_diff.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_add.cu b/test/test_cuda_u32_add.cu index c6ea297..4d6a836 100644 --- a/test/test_cuda_u32_add.cu +++ b/test/test_cuda_u32_add.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_add_error.cu b/test/test_cuda_u32_add_error.cu index 2dd6ec4..2af8a2f 100644 --- a/test/test_cuda_u32_add_error.cu +++ b/test/test_cuda_u32_add_error.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u32_bit_ceil.cu b/test/test_cuda_u32_bit_ceil.cu index 396ec43..0cf574d 100644 --- a/test/test_cuda_u32_bit_ceil.cu +++ b/test/test_cuda_u32_bit_ceil.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_bit_floor.cu b/test/test_cuda_u32_bit_floor.cu index bb9c3e7..0b2a94d 100644 --- a/test/test_cuda_u32_bit_floor.cu +++ b/test/test_cuda_u32_bit_floor.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_bit_width.cu b/test/test_cuda_u32_bit_width.cu index 0ef8388..37ae365 100644 --- a/test/test_cuda_u32_bit_width.cu +++ b/test/test_cuda_u32_bit_width.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_bitswap.cu b/test/test_cuda_u32_bitswap.cu index 3719071..e623f59 100644 --- a/test/test_cuda_u32_bitswap.cu +++ b/test/test_cuda_u32_bitswap.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_byteswap.cu b/test/test_cuda_u32_byteswap.cu index 002fe5e..3931a1f 100644 --- a/test/test_cuda_u32_byteswap.cu +++ b/test/test_cuda_u32_byteswap.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_charconv.cu b/test/test_cuda_u32_charconv.cu index 92084f5..4c4c4d0 100644 --- a/test/test_cuda_u32_charconv.cu +++ b/test/test_cuda_u32_charconv.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_charconv_all_bases.cu b/test/test_cuda_u32_charconv_all_bases.cu index 3f04477..78deac5 100644 --- a/test/test_cuda_u32_charconv_all_bases.cu +++ b/test/test_cuda_u32_charconv_all_bases.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_countl_one.cu b/test/test_cuda_u32_countl_one.cu index 6927f21..78122b2 100644 --- a/test/test_cuda_u32_countl_one.cu +++ b/test/test_cuda_u32_countl_one.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_countl_zero.cu b/test/test_cuda_u32_countl_zero.cu index 474097c..730a5b5 100644 --- a/test/test_cuda_u32_countl_zero.cu +++ b/test/test_cuda_u32_countl_zero.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_countr_one.cu b/test/test_cuda_u32_countr_one.cu index 01c6b8f..11bee24 100644 --- a/test/test_cuda_u32_countr_one.cu +++ b/test/test_cuda_u32_countr_one.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_countr_zero.cu b/test/test_cuda_u32_countr_zero.cu index c1cd361..fe5dc19 100644 --- a/test/test_cuda_u32_countr_zero.cu +++ b/test/test_cuda_u32_countr_zero.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_div.cu b/test/test_cuda_u32_div.cu index 92c7914..6ef31ac 100644 --- a/test/test_cuda_u32_div.cu +++ b/test/test_cuda_u32_div.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_div_ceil.cu b/test/test_cuda_u32_div_ceil.cu index 9cad1a5..0b7e44f 100644 --- a/test/test_cuda_u32_div_ceil.cu +++ b/test/test_cuda_u32_div_ceil.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_div_error.cu b/test/test_cuda_u32_div_error.cu index 423c610..c16a38f 100644 --- a/test/test_cuda_u32_div_error.cu +++ b/test/test_cuda_u32_div_error.cu @@ -8,7 +8,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u32_from_be.cu b/test/test_cuda_u32_from_be.cu index a57e033..d352859 100644 --- a/test/test_cuda_u32_from_be.cu +++ b/test/test_cuda_u32_from_be.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_from_le.cu b/test/test_cuda_u32_from_le.cu index eb3a1f6..8a471a5 100644 --- a/test/test_cuda_u32_from_le.cu +++ b/test/test_cuda_u32_from_le.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_gcd.cu b/test/test_cuda_u32_gcd.cu index 3f1557f..4a3c828 100644 --- a/test/test_cuda_u32_gcd.cu +++ b/test/test_cuda_u32_gcd.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_has_single_bit.cu b/test/test_cuda_u32_has_single_bit.cu index 25b5456..cc6d25e 100644 --- a/test/test_cuda_u32_has_single_bit.cu +++ b/test/test_cuda_u32_has_single_bit.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_ilog.cu b/test/test_cuda_u32_ilog.cu index d4d687a..b9182d1 100644 --- a/test/test_cuda_u32_ilog.cu +++ b/test/test_cuda_u32_ilog.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_ilog10.cu b/test/test_cuda_u32_ilog10.cu index 4d4c865..c4785c9 100644 --- a/test/test_cuda_u32_ilog10.cu +++ b/test/test_cuda_u32_ilog10.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_ilog2.cu b/test/test_cuda_u32_ilog2.cu index 1732664..0bca342 100644 --- a/test/test_cuda_u32_ilog2.cu +++ b/test/test_cuda_u32_ilog2.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_ipow.cu b/test/test_cuda_u32_ipow.cu index 1e41197..b2e157a 100644 --- a/test/test_cuda_u32_ipow.cu +++ b/test/test_cuda_u32_ipow.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_is_power_10.cu b/test/test_cuda_u32_is_power_10.cu index 1fd4d0f..0512be2 100644 --- a/test/test_cuda_u32_is_power_10.cu +++ b/test/test_cuda_u32_is_power_10.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_is_power_2.cu b/test/test_cuda_u32_is_power_2.cu index 6ac5b2d..092b8b5 100644 --- a/test/test_cuda_u32_is_power_2.cu +++ b/test/test_cuda_u32_is_power_2.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_isqrt.cu b/test/test_cuda_u32_isqrt.cu index 3b87830..44ceb25 100644 --- a/test/test_cuda_u32_isqrt.cu +++ b/test/test_cuda_u32_isqrt.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_lcm.cu b/test/test_cuda_u32_lcm.cu index bd77d96..2f1dc09 100644 --- a/test/test_cuda_u32_lcm.cu +++ b/test/test_cuda_u32_lcm.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_midpoint.cu b/test/test_cuda_u32_midpoint.cu index 47901a4..0b6424b 100644 --- a/test/test_cuda_u32_midpoint.cu +++ b/test/test_cuda_u32_midpoint.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_mod.cu b/test/test_cuda_u32_mod.cu index bc3f447..4497132 100644 --- a/test/test_cuda_u32_mod.cu +++ b/test/test_cuda_u32_mod.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_mod_error.cu b/test/test_cuda_u32_mod_error.cu index 9ebaa0c..09f3c84 100644 --- a/test/test_cuda_u32_mod_error.cu +++ b/test/test_cuda_u32_mod_error.cu @@ -8,7 +8,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u32_mul.cu b/test/test_cuda_u32_mul.cu index 2947ac5..44526e8 100644 --- a/test/test_cuda_u32_mul.cu +++ b/test/test_cuda_u32_mul.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_mul_error.cu b/test/test_cuda_u32_mul_error.cu index 58a4596..9096fb0 100644 --- a/test/test_cuda_u32_mul_error.cu +++ b/test/test_cuda_u32_mul_error.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u32_next_multiple_of.cu b/test/test_cuda_u32_next_multiple_of.cu index fa16fd1..7c74e01 100644 --- a/test/test_cuda_u32_next_multiple_of.cu +++ b/test/test_cuda_u32_next_multiple_of.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_popcount.cu b/test/test_cuda_u32_popcount.cu index a9eed9c..dd18fe6 100644 --- a/test/test_cuda_u32_popcount.cu +++ b/test/test_cuda_u32_popcount.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_remove_trailing_zeros.cu b/test/test_cuda_u32_remove_trailing_zeros.cu index 3b52b97..115c361 100644 --- a/test/test_cuda_u32_remove_trailing_zeros.cu +++ b/test/test_cuda_u32_remove_trailing_zeros.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_rotl.cu b/test/test_cuda_u32_rotl.cu index aada247..e82d70c 100644 --- a/test/test_cuda_u32_rotl.cu +++ b/test/test_cuda_u32_rotl.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_rotr.cu b/test/test_cuda_u32_rotr.cu index a38c382..3d4036f 100644 --- a/test/test_cuda_u32_rotr.cu +++ b/test/test_cuda_u32_rotr.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_sub.cu b/test/test_cuda_u32_sub.cu index d5d4b6e..129e498 100644 --- a/test/test_cuda_u32_sub.cu +++ b/test/test_cuda_u32_sub.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_sub_error.cu b/test/test_cuda_u32_sub_error.cu index 40f4c0b..bc1a696 100644 --- a/test/test_cuda_u32_sub_error.cu +++ b/test/test_cuda_u32_sub_error.cu @@ -8,7 +8,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u32_to_be.cu b/test/test_cuda_u32_to_be.cu index 16b7861..337bc9c 100644 --- a/test/test_cuda_u32_to_be.cu +++ b/test/test_cuda_u32_to_be.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u32_to_le.cu b/test/test_cuda_u32_to_le.cu index b685cde..a6a73bf 100644 --- a/test/test_cuda_u32_to_le.cu +++ b/test/test_cuda_u32_to_le.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_abs_diff.cu b/test/test_cuda_u64_abs_diff.cu index 1fb8658..4fe66d9 100644 --- a/test/test_cuda_u64_abs_diff.cu +++ b/test/test_cuda_u64_abs_diff.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_add.cu b/test/test_cuda_u64_add.cu index 7cd2ebe..10d8cc7 100644 --- a/test/test_cuda_u64_add.cu +++ b/test/test_cuda_u64_add.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_add_error.cu b/test/test_cuda_u64_add_error.cu index 45afb1b..b1d3052 100644 --- a/test/test_cuda_u64_add_error.cu +++ b/test/test_cuda_u64_add_error.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u64_bit_ceil.cu b/test/test_cuda_u64_bit_ceil.cu index 6b546f6..a54e8fc 100644 --- a/test/test_cuda_u64_bit_ceil.cu +++ b/test/test_cuda_u64_bit_ceil.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_bit_floor.cu b/test/test_cuda_u64_bit_floor.cu index c5c5344..cac52a2 100644 --- a/test/test_cuda_u64_bit_floor.cu +++ b/test/test_cuda_u64_bit_floor.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_bit_width.cu b/test/test_cuda_u64_bit_width.cu index 333df34..d9bd633 100644 --- a/test/test_cuda_u64_bit_width.cu +++ b/test/test_cuda_u64_bit_width.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_bitswap.cu b/test/test_cuda_u64_bitswap.cu index 22860ae..83cb272 100644 --- a/test/test_cuda_u64_bitswap.cu +++ b/test/test_cuda_u64_bitswap.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_byteswap.cu b/test/test_cuda_u64_byteswap.cu index 88b81e2..58554b0 100644 --- a/test/test_cuda_u64_byteswap.cu +++ b/test/test_cuda_u64_byteswap.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_charconv.cu b/test/test_cuda_u64_charconv.cu index b20ff71..0a472bb 100644 --- a/test/test_cuda_u64_charconv.cu +++ b/test/test_cuda_u64_charconv.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_charconv_all_bases.cu b/test/test_cuda_u64_charconv_all_bases.cu index b05fadc..4875c4d 100644 --- a/test/test_cuda_u64_charconv_all_bases.cu +++ b/test/test_cuda_u64_charconv_all_bases.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_countl_one.cu b/test/test_cuda_u64_countl_one.cu index 2fa797c..7c83a36 100644 --- a/test/test_cuda_u64_countl_one.cu +++ b/test/test_cuda_u64_countl_one.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_countl_zero.cu b/test/test_cuda_u64_countl_zero.cu index f5ee174..16fcdb3 100644 --- a/test/test_cuda_u64_countl_zero.cu +++ b/test/test_cuda_u64_countl_zero.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_countr_one.cu b/test/test_cuda_u64_countr_one.cu index b6f2dd1..d923a06 100644 --- a/test/test_cuda_u64_countr_one.cu +++ b/test/test_cuda_u64_countr_one.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_countr_zero.cu b/test/test_cuda_u64_countr_zero.cu index de18f07..f6be2db 100644 --- a/test/test_cuda_u64_countr_zero.cu +++ b/test/test_cuda_u64_countr_zero.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_div.cu b/test/test_cuda_u64_div.cu index ad4adc6..6353ea4 100644 --- a/test/test_cuda_u64_div.cu +++ b/test/test_cuda_u64_div.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_div_ceil.cu b/test/test_cuda_u64_div_ceil.cu index 69facdf..c3bc246 100644 --- a/test/test_cuda_u64_div_ceil.cu +++ b/test/test_cuda_u64_div_ceil.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_div_error.cu b/test/test_cuda_u64_div_error.cu index 27f8697..b43a236 100644 --- a/test/test_cuda_u64_div_error.cu +++ b/test/test_cuda_u64_div_error.cu @@ -8,7 +8,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u64_from_be.cu b/test/test_cuda_u64_from_be.cu index becb752..a5b6406 100644 --- a/test/test_cuda_u64_from_be.cu +++ b/test/test_cuda_u64_from_be.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_from_le.cu b/test/test_cuda_u64_from_le.cu index 96846e0..0c6940a 100644 --- a/test/test_cuda_u64_from_le.cu +++ b/test/test_cuda_u64_from_le.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_gcd.cu b/test/test_cuda_u64_gcd.cu index 14b93ee..c84c157 100644 --- a/test/test_cuda_u64_gcd.cu +++ b/test/test_cuda_u64_gcd.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_has_single_bit.cu b/test/test_cuda_u64_has_single_bit.cu index febd906..6258d9e 100644 --- a/test/test_cuda_u64_has_single_bit.cu +++ b/test/test_cuda_u64_has_single_bit.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_ilog.cu b/test/test_cuda_u64_ilog.cu index 8aec005..fcaf66a 100644 --- a/test/test_cuda_u64_ilog.cu +++ b/test/test_cuda_u64_ilog.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_ilog10.cu b/test/test_cuda_u64_ilog10.cu index c26fb8d..4e9c1e0 100644 --- a/test/test_cuda_u64_ilog10.cu +++ b/test/test_cuda_u64_ilog10.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_ilog2.cu b/test/test_cuda_u64_ilog2.cu index 111a156..c682e71 100644 --- a/test/test_cuda_u64_ilog2.cu +++ b/test/test_cuda_u64_ilog2.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_ipow.cu b/test/test_cuda_u64_ipow.cu index 40182ad..82070f9 100644 --- a/test/test_cuda_u64_ipow.cu +++ b/test/test_cuda_u64_ipow.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_is_power_10.cu b/test/test_cuda_u64_is_power_10.cu index 09e883e..08883bb 100644 --- a/test/test_cuda_u64_is_power_10.cu +++ b/test/test_cuda_u64_is_power_10.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_is_power_2.cu b/test/test_cuda_u64_is_power_2.cu index d637dfa..8c01d57 100644 --- a/test/test_cuda_u64_is_power_2.cu +++ b/test/test_cuda_u64_is_power_2.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_isqrt.cu b/test/test_cuda_u64_isqrt.cu index 03d3281..0059ad8 100644 --- a/test/test_cuda_u64_isqrt.cu +++ b/test/test_cuda_u64_isqrt.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_lcm.cu b/test/test_cuda_u64_lcm.cu index a9b8902..fed2a98 100644 --- a/test/test_cuda_u64_lcm.cu +++ b/test/test_cuda_u64_lcm.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_midpoint.cu b/test/test_cuda_u64_midpoint.cu index 5b98709..f4717c1 100644 --- a/test/test_cuda_u64_midpoint.cu +++ b/test/test_cuda_u64_midpoint.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_mod.cu b/test/test_cuda_u64_mod.cu index 5f937a5..db09f3e 100644 --- a/test/test_cuda_u64_mod.cu +++ b/test/test_cuda_u64_mod.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_mod_error.cu b/test/test_cuda_u64_mod_error.cu index 2f877f9..29081b3 100644 --- a/test/test_cuda_u64_mod_error.cu +++ b/test/test_cuda_u64_mod_error.cu @@ -8,7 +8,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u64_mul.cu b/test/test_cuda_u64_mul.cu index 6a7c86d..6e73530 100644 --- a/test/test_cuda_u64_mul.cu +++ b/test/test_cuda_u64_mul.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_mul_error.cu b/test/test_cuda_u64_mul_error.cu index aee86b4..6591e80 100644 --- a/test/test_cuda_u64_mul_error.cu +++ b/test/test_cuda_u64_mul_error.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u64_next_multiple_of.cu b/test/test_cuda_u64_next_multiple_of.cu index b18eadc..c2711f5 100644 --- a/test/test_cuda_u64_next_multiple_of.cu +++ b/test/test_cuda_u64_next_multiple_of.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_popcount.cu b/test/test_cuda_u64_popcount.cu index a9d883b..d313b87 100644 --- a/test/test_cuda_u64_popcount.cu +++ b/test/test_cuda_u64_popcount.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_remove_trailing_zeros.cu b/test/test_cuda_u64_remove_trailing_zeros.cu index 3ca4ffd..8c27bb6 100644 --- a/test/test_cuda_u64_remove_trailing_zeros.cu +++ b/test/test_cuda_u64_remove_trailing_zeros.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_rotl.cu b/test/test_cuda_u64_rotl.cu index ac84552..1000dba 100644 --- a/test/test_cuda_u64_rotl.cu +++ b/test/test_cuda_u64_rotl.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_rotr.cu b/test/test_cuda_u64_rotr.cu index 017cb12..3253f3a 100644 --- a/test/test_cuda_u64_rotr.cu +++ b/test/test_cuda_u64_rotr.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_sub.cu b/test/test_cuda_u64_sub.cu index 5e235a1..2700515 100644 --- a/test/test_cuda_u64_sub.cu +++ b/test/test_cuda_u64_sub.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_sub_error.cu b/test/test_cuda_u64_sub_error.cu index fc69b82..826d71c 100644 --- a/test/test_cuda_u64_sub_error.cu +++ b/test/test_cuda_u64_sub_error.cu @@ -8,7 +8,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u64_to_be.cu b/test/test_cuda_u64_to_be.cu index 47bed5b..124fda6 100644 --- a/test/test_cuda_u64_to_be.cu +++ b/test/test_cuda_u64_to_be.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u64_to_le.cu b/test/test_cuda_u64_to_le.cu index d5757d1..d71d93f 100644 --- a/test/test_cuda_u64_to_le.cu +++ b/test/test_cuda_u64_to_le.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_abs_diff.cu b/test/test_cuda_u8_abs_diff.cu index 2ae8bfc..89cd086 100644 --- a/test/test_cuda_u8_abs_diff.cu +++ b/test/test_cuda_u8_abs_diff.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_add.cu b/test/test_cuda_u8_add.cu index 77df882..d00aacd 100644 --- a/test/test_cuda_u8_add.cu +++ b/test/test_cuda_u8_add.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_add_error.cu b/test/test_cuda_u8_add_error.cu index 9fb4a6b..18d150c 100644 --- a/test/test_cuda_u8_add_error.cu +++ b/test/test_cuda_u8_add_error.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u8_bit_ceil.cu b/test/test_cuda_u8_bit_ceil.cu index 0eaeb22..d3ef354 100644 --- a/test/test_cuda_u8_bit_ceil.cu +++ b/test/test_cuda_u8_bit_ceil.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_bit_floor.cu b/test/test_cuda_u8_bit_floor.cu index c3d4c0c..04e282f 100644 --- a/test/test_cuda_u8_bit_floor.cu +++ b/test/test_cuda_u8_bit_floor.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_bit_width.cu b/test/test_cuda_u8_bit_width.cu index 4edc3ca..4bda9c6 100644 --- a/test/test_cuda_u8_bit_width.cu +++ b/test/test_cuda_u8_bit_width.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_bitswap.cu b/test/test_cuda_u8_bitswap.cu index 726a52d..0ec19ff 100644 --- a/test/test_cuda_u8_bitswap.cu +++ b/test/test_cuda_u8_bitswap.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_byteswap.cu b/test/test_cuda_u8_byteswap.cu index 59921b9..4dc5af8 100644 --- a/test/test_cuda_u8_byteswap.cu +++ b/test/test_cuda_u8_byteswap.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_charconv.cu b/test/test_cuda_u8_charconv.cu index b1756d0..47bdcf9 100644 --- a/test/test_cuda_u8_charconv.cu +++ b/test/test_cuda_u8_charconv.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_charconv_all_bases.cu b/test/test_cuda_u8_charconv_all_bases.cu index 9a08503..c0315be 100644 --- a/test/test_cuda_u8_charconv_all_bases.cu +++ b/test/test_cuda_u8_charconv_all_bases.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_countl_one.cu b/test/test_cuda_u8_countl_one.cu index 0a535ba..162bd7a 100644 --- a/test/test_cuda_u8_countl_one.cu +++ b/test/test_cuda_u8_countl_one.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_countl_zero.cu b/test/test_cuda_u8_countl_zero.cu index 6643b7b..749206d 100644 --- a/test/test_cuda_u8_countl_zero.cu +++ b/test/test_cuda_u8_countl_zero.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_countr_one.cu b/test/test_cuda_u8_countr_one.cu index 73a7e5b..1529cfd 100644 --- a/test/test_cuda_u8_countr_one.cu +++ b/test/test_cuda_u8_countr_one.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_countr_zero.cu b/test/test_cuda_u8_countr_zero.cu index e722b4c..3db60fa 100644 --- a/test/test_cuda_u8_countr_zero.cu +++ b/test/test_cuda_u8_countr_zero.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_div.cu b/test/test_cuda_u8_div.cu index 9bfd9b5..08fe105 100644 --- a/test/test_cuda_u8_div.cu +++ b/test/test_cuda_u8_div.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_div_ceil.cu b/test/test_cuda_u8_div_ceil.cu index f2f8cbb..4304b89 100644 --- a/test/test_cuda_u8_div_ceil.cu +++ b/test/test_cuda_u8_div_ceil.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_div_error.cu b/test/test_cuda_u8_div_error.cu index 5cddd3b..f3f6b85 100644 --- a/test/test_cuda_u8_div_error.cu +++ b/test/test_cuda_u8_div_error.cu @@ -8,7 +8,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u8_from_be.cu b/test/test_cuda_u8_from_be.cu index 9e0e941..630be25 100644 --- a/test/test_cuda_u8_from_be.cu +++ b/test/test_cuda_u8_from_be.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_from_le.cu b/test/test_cuda_u8_from_le.cu index 7676cb9..7e65638 100644 --- a/test/test_cuda_u8_from_le.cu +++ b/test/test_cuda_u8_from_le.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_gcd.cu b/test/test_cuda_u8_gcd.cu index 3c9e888..4c73d88 100644 --- a/test/test_cuda_u8_gcd.cu +++ b/test/test_cuda_u8_gcd.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_has_single_bit.cu b/test/test_cuda_u8_has_single_bit.cu index f716108..189c36e 100644 --- a/test/test_cuda_u8_has_single_bit.cu +++ b/test/test_cuda_u8_has_single_bit.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_ilog.cu b/test/test_cuda_u8_ilog.cu index c47ce6b..c7ff6eb 100644 --- a/test/test_cuda_u8_ilog.cu +++ b/test/test_cuda_u8_ilog.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_ilog10.cu b/test/test_cuda_u8_ilog10.cu index 47b296a..9a06420 100644 --- a/test/test_cuda_u8_ilog10.cu +++ b/test/test_cuda_u8_ilog10.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_ilog2.cu b/test/test_cuda_u8_ilog2.cu index 49926f9..02daa3f 100644 --- a/test/test_cuda_u8_ilog2.cu +++ b/test/test_cuda_u8_ilog2.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_ipow.cu b/test/test_cuda_u8_ipow.cu index be222da..a8859d4 100644 --- a/test/test_cuda_u8_ipow.cu +++ b/test/test_cuda_u8_ipow.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_is_power_10.cu b/test/test_cuda_u8_is_power_10.cu index 84ea12b..ea51e19 100644 --- a/test/test_cuda_u8_is_power_10.cu +++ b/test/test_cuda_u8_is_power_10.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_is_power_2.cu b/test/test_cuda_u8_is_power_2.cu index a4551ed..f34a783 100644 --- a/test/test_cuda_u8_is_power_2.cu +++ b/test/test_cuda_u8_is_power_2.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_isqrt.cu b/test/test_cuda_u8_isqrt.cu index cc39059..3d5f4c2 100644 --- a/test/test_cuda_u8_isqrt.cu +++ b/test/test_cuda_u8_isqrt.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_lcm.cu b/test/test_cuda_u8_lcm.cu index 1b895c1..15b0f71 100644 --- a/test/test_cuda_u8_lcm.cu +++ b/test/test_cuda_u8_lcm.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_midpoint.cu b/test/test_cuda_u8_midpoint.cu index b4521b3..829d031 100644 --- a/test/test_cuda_u8_midpoint.cu +++ b/test/test_cuda_u8_midpoint.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_mod.cu b/test/test_cuda_u8_mod.cu index 55f11b6..e4015d6 100644 --- a/test/test_cuda_u8_mod.cu +++ b/test/test_cuda_u8_mod.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_mod_error.cu b/test/test_cuda_u8_mod_error.cu index bc9d668..2af089f 100644 --- a/test/test_cuda_u8_mod_error.cu +++ b/test/test_cuda_u8_mod_error.cu @@ -8,7 +8,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u8_mul.cu b/test/test_cuda_u8_mul.cu index 8ae9cc4..e47aa85 100644 --- a/test/test_cuda_u8_mul.cu +++ b/test/test_cuda_u8_mul.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_mul_error.cu b/test/test_cuda_u8_mul_error.cu index 79101bc..4d1db43 100644 --- a/test/test_cuda_u8_mul_error.cu +++ b/test/test_cuda_u8_mul_error.cu @@ -9,7 +9,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u8_next_multiple_of.cu b/test/test_cuda_u8_next_multiple_of.cu index b676786..fc4f3b3 100644 --- a/test/test_cuda_u8_next_multiple_of.cu +++ b/test/test_cuda_u8_next_multiple_of.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_popcount.cu b/test/test_cuda_u8_popcount.cu index 6c5ba4f..fc3f0f6 100644 --- a/test/test_cuda_u8_popcount.cu +++ b/test/test_cuda_u8_popcount.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_remove_trailing_zeros.cu b/test/test_cuda_u8_remove_trailing_zeros.cu index 7b4b9f2..e28a7f9 100644 --- a/test/test_cuda_u8_remove_trailing_zeros.cu +++ b/test/test_cuda_u8_remove_trailing_zeros.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_rotl.cu b/test/test_cuda_u8_rotl.cu index 28e06a4..38953e9 100644 --- a/test/test_cuda_u8_rotl.cu +++ b/test/test_cuda_u8_rotl.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_rotr.cu b/test/test_cuda_u8_rotr.cu index b50db99..90e4c42 100644 --- a/test/test_cuda_u8_rotr.cu +++ b/test/test_cuda_u8_rotr.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_sub.cu b/test/test_cuda_u8_sub.cu index 4036564..b16c9bc 100644 --- a/test/test_cuda_u8_sub.cu +++ b/test/test_cuda_u8_sub.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_sub_error.cu b/test/test_cuda_u8_sub_error.cu index 045f91b..6ac0ec1 100644 --- a/test/test_cuda_u8_sub_error.cu +++ b/test/test_cuda_u8_sub_error.cu @@ -8,7 +8,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include diff --git a/test/test_cuda_u8_to_be.cu b/test/test_cuda_u8_to_be.cu index 7b184c9..c900255 100644 --- a/test/test_cuda_u8_to_be.cu +++ b/test/test_cuda_u8_to_be.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_cuda_u8_to_le.cu b/test/test_cuda_u8_to_le.cu index 2de11d0..c8c4211 100644 --- a/test/test_cuda_u8_to_le.cu +++ b/test/test_cuda_u8_to_le.cu @@ -10,7 +10,7 @@ #include #include #include -#include +#include #include "cuda_managed_ptr.hpp" #include "stopwatch.hpp" diff --git a/test/test_device_error_context_reuse_sycl.cpp b/test/test_device_error_context_reuse_sycl.cpp new file mode 100644 index 0000000..ddddf52 --- /dev/null +++ b/test/test_device_error_context_reuse_sycl.cpp @@ -0,0 +1,86 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// After synchronize() throws a captured device error, the context clears the +// device buffer and must be immediately reusable for further launches (the +// SYCL analog of the CUDA untrapped-mode reuse test). + +#include "sycl_test.hpp" +#include + +int main() +{ + using boost::safe_numbers::u32; + + sycl::queue q; + boost::safe_numbers::device_error_context ctx {q}; + + constexpr int n {64}; + auto* big {sycl::malloc_shared(n, q)}; + auto* out {sycl::malloc_shared(n, q)}; + for (int i {0}; i < n; ++i) + { + big[i] = u32{(std::numeric_limits::max)()}; + } + + // First launch overflows and must throw on synchronize + q.parallel_for(sycl::range<1>(n), [=](sycl::id<1> idx) + { + const int i {static_cast(idx[0])}; + out[i] = big[i] + big[i]; + }); + + bool caught {false}; + try + { + ctx.synchronize(); + } + catch (const std::overflow_error&) + { + caught = true; + } + + if (!caught) + { + std::cerr << "FAILED: no overflow_error from the first launch\n"; + sycl::free(big, q); + sycl::free(out, q); + return EXIT_FAILURE; + } + + // Second launch on the same context must run cleanly and synchronize + // without throwing + q.parallel_for(sycl::range<1>(n), [=](sycl::id<1> idx) + { + const int i {static_cast(idx[0])}; + out[i] = u32{1U} + u32{2U}; + }); + + try + { + ctx.synchronize(); + } + catch (const std::exception& e) + { + std::cerr << "FAILED: context not reusable: " << e.what() << "\n"; + sycl::free(big, q); + sycl::free(out, q); + return EXIT_FAILURE; + } + + int failures {0}; + for (int i {0}; i < n; ++i) + { + if (!(out[i] == u32{3U})) + { + ++failures; + } + } + + sycl::free(big, q); + sycl::free(out, q); + + std::cout << "context reuse: " << (failures == 0 ? "PASSED" : "FAILED") << "\n"; + return failures == 0 ? EXIT_SUCCESS : EXIT_FAILURE; +} diff --git a/test/test_device_error_mapping.cpp b/test/test_device_error_mapping.cpp index 2b52044..047fd96 100644 --- a/test/test_device_error_mapping.cpp +++ b/test/test_device_error_mapping.cpp @@ -8,7 +8,7 @@ // own category rather than being folded into domain_error, so that a device // (or CUDA host) invalid_argument is rethrown as std::invalid_argument. -#include +#include #include using boost::safe_numbers::detail::exception_type; diff --git a/test/test_device_error_mapping_sycl.cpp b/test/test_device_error_mapping_sycl.cpp new file mode 100644 index 0000000..e7096ac --- /dev/null +++ b/test/test_device_error_mapping_sycl.cpp @@ -0,0 +1,52 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// Reporting through the device funnel directly must map every +// exception_type enumerator to the matching std exception on the host +// (the SYCL analog of the CUDA invalid_argument round-trip test). + +#include "sycl_test.hpp" + +template +int run_direct_report(const char* label, const boost::safe_numbers::detail::exception_type exc) +{ + sycl::queue q; + boost::safe_numbers::device_error_context ctx {q}; + + q.single_task([=]() + { + boost::safe_numbers::detail::report_device_error(exc, "mapping_test", 1, "direct report mapping"); + }); + + try + { + ctx.synchronize(); + } + catch (const Expected& e) + { + std::cout << label << ": PASSED (" << e.what() << ")\n"; + return EXIT_SUCCESS; + } + catch (const std::exception& e) + { + std::cerr << label << ": FAILED - wrong exception type: " << e.what() << "\n"; + return EXIT_FAILURE; + } + + std::cerr << label << ": FAILED - no exception thrown\n"; + return EXIT_FAILURE; +} + +int main() +{ + using boost::safe_numbers::detail::exception_type; + + int rc {EXIT_SUCCESS}; + rc |= run_direct_report("domain_error mapping", exception_type::domain_error); + rc |= run_direct_report("overflow_error mapping", exception_type::overflow); + rc |= run_direct_report("underflow_error mapping", exception_type::underflow); + rc |= run_direct_report("invalid_argument mapping", exception_type::invalid_argument); + rc |= run_direct_report("unknown mapping", exception_type::unknown); + return rc; +} diff --git a/test/test_device_error_single_context_sycl.cpp b/test/test_device_error_single_context_sycl.cpp new file mode 100644 index 0000000..8a1b162 --- /dev/null +++ b/test/test_device_error_single_context_sycl.cpp @@ -0,0 +1,51 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// Only one device_error_context may exist at a time (the device buffer is a +// single device_global shared by all contexts). A second simultaneous +// context must throw std::logic_error; sequential contexts are fine. + +#include "sycl_test.hpp" + +int main() +{ + sycl::queue q; + + int rc {EXIT_SUCCESS}; + + { + boost::safe_numbers::device_error_context ctx {q}; + + bool caught {false}; + try + { + boost::safe_numbers::device_error_context second {q}; + } + catch (const std::logic_error&) + { + caught = true; + } + + if (!caught) + { + std::cerr << "FAILED: second simultaneous context did not throw std::logic_error\n"; + rc = EXIT_FAILURE; + } + } + + // After the first context is destroyed a new one may be created + try + { + boost::safe_numbers::device_error_context ctx {q}; + ctx.synchronize(); + } + catch (const std::exception& e) + { + std::cerr << "FAILED: sequential context construction threw: " << e.what() << "\n"; + rc = EXIT_FAILURE; + } + + std::cout << "single context rule: " << (rc == EXIT_SUCCESS ? "PASSED" : "FAILED") << "\n"; + return rc; +} diff --git a/test/test_device_error_workitem_id_sycl.cpp b/test/test_device_error_workitem_id_sycl.cpp new file mode 100644 index 0000000..38c5c8d --- /dev/null +++ b/test/test_device_error_workitem_id_sycl.cpp @@ -0,0 +1,90 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// When exactly one element triggers a device error, the recorded work-item +// id must identify it, so the failing input can be narrowed down exactly as +// with the CUDA backend. Checked for both a plain range launch and an +// nd_range launch. + +#include "sycl_test.hpp" +#include +#include + +using boost::safe_numbers::u32; + +constexpr int n {4096}; + +template +int expect_thread_in_message(const char* label, const int bad_index, Launch launch) +{ + sycl::queue q; + boost::safe_numbers::device_error_context ctx {q}; + + auto* a {sycl::malloc_shared(n, q)}; + auto* out {sycl::malloc_shared(n, q)}; + for (int i {0}; i < n; ++i) + { + // Only a[bad_index] overflows when doubled + a[i] = (i == bad_index) ? u32{(std::numeric_limits::max)()} : u32{1U}; + } + + launch(q, a, out); + + int rc {EXIT_FAILURE}; + try + { + ctx.synchronize(); + std::cerr << label << ": FAILED - no exception thrown\n"; + } + catch (const std::overflow_error& e) + { + const std::string what {e.what()}; + const std::string expected {"thread " + std::to_string(bad_index)}; + if (what.find(expected) != std::string::npos) + { + std::cout << label << ": PASSED (" << what << ")\n"; + rc = EXIT_SUCCESS; + } + else + { + std::cerr << label << ": FAILED - expected '" << expected << "' in: " << what << "\n"; + } + } + catch (const std::exception& e) + { + std::cerr << label << ": FAILED - wrong exception type: " << e.what() << "\n"; + } + + sycl::free(a, q); + sycl::free(out, q); + return rc; +} + +int main() +{ + int rc {EXIT_SUCCESS}; + + rc |= expect_thread_in_message("plain range work-item id", 1234, + [](sycl::queue& q, const u32* a, u32* out) + { + q.parallel_for(sycl::range<1>(n), [=](sycl::id<1> idx) + { + const int i {static_cast(idx[0])}; + out[i] = a[i] + a[i]; + }); + }); + + rc |= expect_thread_in_message("nd_range work-item id", 777, + [](sycl::queue& q, const u32* a, u32* out) + { + q.parallel_for(sycl::nd_range<1>(sycl::range<1>(n), sycl::range<1>(64)), + [=](sycl::nd_item<1> it) + { + const int i {static_cast(it.get_global_linear_id())}; + out[i] = a[i] + a[i]; + }); + }); + + return rc; +} diff --git a/test/test_div_ceil_sycl.cpp b/test/test_div_ceil_sycl.cpp new file mode 100644 index 0000000..36a3dec --- /dev/null +++ b/test/test_div_ceil_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_binary_all_unsigned("div_ceil", [](auto a, auto b) { return boost::safe_numbers::div_ceil(a, b); }); +} diff --git a/test/test_float_add_error_sycl.cpp b/test/test_float_add_error_sycl.cpp new file mode 100644 index 0000000..1458c17 --- /dev/null +++ b/test/test_float_add_error_sycl.cpp @@ -0,0 +1,29 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// A float addition that saturates to +infinity must be reported from the device and surface on the host as +// std::overflow_error via device_error_context::synchronize(). + +#include "sycl_test.hpp" +#include + +template +int run(const char* label) +{ + using basis_type = typename T::basis_type; + + return sn_sycl_test::expect_binary_throw(label, + T{(std::numeric_limits::max)()}, T{(std::numeric_limits::max)()}, + [](auto x, auto y) { return x + y; }); +} + +int main() +{ + using namespace boost::safe_numbers; + + int rc {EXIT_SUCCESS}; + rc |= run("f32 add"); + rc |= run("f64 add"); + return rc; +} diff --git a/test/test_float_add_sycl.cpp b/test/test_float_add_sycl.cpp new file mode 100644 index 0000000..b2c893c --- /dev/null +++ b/test/test_float_add_sycl.cpp @@ -0,0 +1,27 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// Addition on the device for f32 and f64 over random finite operands, +// verified exactly against a host recomputation. + +#include "sycl_test.hpp" + +template +int run(const char* label) +{ + return sn_sycl_test::run_binary_2(label, + [](std::mt19937_64& rng) { return sn_sycl_test::draw_real(rng, -1e18, 1e18); }, + [](std::mt19937_64& rng) { return sn_sycl_test::draw_real(rng, -1e18, 1e18); }, + [](auto a, auto b) { return a + b; }); +} + +int main() +{ + using namespace boost::safe_numbers; + + int rc {EXIT_SUCCESS}; + rc |= run("f32 add"); + rc |= run("f64 add"); + return rc; +} diff --git a/test/test_float_div_error_sycl.cpp b/test/test_float_div_error_sycl.cpp new file mode 100644 index 0000000..ed34c5d --- /dev/null +++ b/test/test_float_div_error_sycl.cpp @@ -0,0 +1,29 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// A float division by zero must be reported from the device and surface on the host as +// std::domain_error via device_error_context::synchronize(). + +#include "sycl_test.hpp" +#include + +template +int run(const char* label) +{ + using basis_type = typename T::basis_type; + + return sn_sycl_test::expect_binary_throw(label, + T{static_cast(1)}, T{static_cast(0)}, + [](auto x, auto y) { return x / y; }); +} + +int main() +{ + using namespace boost::safe_numbers; + + int rc {EXIT_SUCCESS}; + rc |= run("f32 div"); + rc |= run("f64 div"); + return rc; +} diff --git a/test/test_float_div_sycl.cpp b/test/test_float_div_sycl.cpp new file mode 100644 index 0000000..26f30df --- /dev/null +++ b/test/test_float_div_sycl.cpp @@ -0,0 +1,27 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// Division on the device for f32 and f64 over random finite operands, +// verified exactly against a host recomputation. + +#include "sycl_test.hpp" + +template +int run(const char* label) +{ + return sn_sycl_test::run_binary_2(label, + [](std::mt19937_64& rng) { return sn_sycl_test::draw_real(rng, -1e15, 1e15); }, + [](std::mt19937_64& rng) { return sn_sycl_test::draw_real(rng, 1, 1e15); }, + [](auto a, auto b) { return a / b; }); +} + +int main() +{ + using namespace boost::safe_numbers; + + int rc {EXIT_SUCCESS}; + rc |= run("f32 div"); + rc |= run("f64 div"); + return rc; +} diff --git a/test/test_float_mul_error_sycl.cpp b/test/test_float_mul_error_sycl.cpp new file mode 100644 index 0000000..d923049 --- /dev/null +++ b/test/test_float_mul_error_sycl.cpp @@ -0,0 +1,29 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// A float multiplication that saturates to +infinity must be reported from the device and surface on the host as +// std::overflow_error via device_error_context::synchronize(). + +#include "sycl_test.hpp" +#include + +template +int run(const char* label) +{ + using basis_type = typename T::basis_type; + + return sn_sycl_test::expect_binary_throw(label, + T{(std::numeric_limits::max)()}, T{static_cast(2)}, + [](auto x, auto y) { return x * y; }); +} + +int main() +{ + using namespace boost::safe_numbers; + + int rc {EXIT_SUCCESS}; + rc |= run("f32 mul"); + rc |= run("f64 mul"); + return rc; +} diff --git a/test/test_float_mul_sycl.cpp b/test/test_float_mul_sycl.cpp new file mode 100644 index 0000000..9e2da8d --- /dev/null +++ b/test/test_float_mul_sycl.cpp @@ -0,0 +1,27 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// Multiplication on the device for f32 and f64 over random finite operands, +// verified exactly against a host recomputation. + +#include "sycl_test.hpp" + +template +int run(const char* label) +{ + return sn_sycl_test::run_binary_2(label, + [](std::mt19937_64& rng) { return sn_sycl_test::draw_real(rng, -1e15, 1e15); }, + [](std::mt19937_64& rng) { return sn_sycl_test::draw_real(rng, -1e15, 1e15); }, + [](auto a, auto b) { return a * b; }); +} + +int main() +{ + using namespace boost::safe_numbers; + + int rc {EXIT_SUCCESS}; + rc |= run("f32 mul"); + rc |= run("f64 mul"); + return rc; +} diff --git a/test/test_float_sub_error_sycl.cpp b/test/test_float_sub_error_sycl.cpp new file mode 100644 index 0000000..967e627 --- /dev/null +++ b/test/test_float_sub_error_sycl.cpp @@ -0,0 +1,29 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// A float subtraction that saturates to -infinity must be reported from the device and surface on the host as +// std::underflow_error via device_error_context::synchronize(). + +#include "sycl_test.hpp" +#include + +template +int run(const char* label) +{ + using basis_type = typename T::basis_type; + + return sn_sycl_test::expect_binary_throw(label, + T{std::numeric_limits::lowest()}, T{(std::numeric_limits::max)()}, + [](auto x, auto y) { return x - y; }); +} + +int main() +{ + using namespace boost::safe_numbers; + + int rc {EXIT_SUCCESS}; + rc |= run("f32 sub"); + rc |= run("f64 sub"); + return rc; +} diff --git a/test/test_float_sub_sycl.cpp b/test/test_float_sub_sycl.cpp new file mode 100644 index 0000000..74152fa --- /dev/null +++ b/test/test_float_sub_sycl.cpp @@ -0,0 +1,27 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// Subtraction on the device for f32 and f64 over random finite operands, +// verified exactly against a host recomputation. + +#include "sycl_test.hpp" + +template +int run(const char* label) +{ + return sn_sycl_test::run_binary_2(label, + [](std::mt19937_64& rng) { return sn_sycl_test::draw_real(rng, -1e18, 1e18); }, + [](std::mt19937_64& rng) { return sn_sycl_test::draw_real(rng, -1e18, 1e18); }, + [](auto a, auto b) { return a - b; }); +} + +int main() +{ + using namespace boost::safe_numbers; + + int rc {EXIT_SUCCESS}; + rc |= run("f32 sub"); + rc |= run("f64 sub"); + return rc; +} diff --git a/test/test_from_be_sycl.cpp b/test/test_from_be_sycl.cpp new file mode 100644 index 0000000..fabae93 --- /dev/null +++ b/test/test_from_be_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("from_be", [](auto a) { return boost::safe_numbers::from_be(a); }); +} diff --git a/test/test_from_le_sycl.cpp b/test/test_from_le_sycl.cpp new file mode 100644 index 0000000..fd92b41 --- /dev/null +++ b/test/test_from_le_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("from_le", [](auto a) { return boost::safe_numbers::from_le(a); }); +} diff --git a/test/test_gcd_sycl.cpp b/test/test_gcd_sycl.cpp new file mode 100644 index 0000000..a559b74 --- /dev/null +++ b/test/test_gcd_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_binary_all_unsigned("gcd", [](auto a, auto b) { return boost::safe_numbers::gcd(a, b); }); +} diff --git a/test/test_has_single_bit_sycl.cpp b/test/test_has_single_bit_sycl.cpp new file mode 100644 index 0000000..a1cfaa5 --- /dev/null +++ b/test/test_has_single_bit_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("has_single_bit", [](auto a) { return static_cast(boost::safe_numbers::has_single_bit(a)); }); +} diff --git a/test/test_ilog10_sycl.cpp b/test/test_ilog10_sycl.cpp new file mode 100644 index 0000000..cf2a278 --- /dev/null +++ b/test/test_ilog10_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("ilog10", [](auto a) { return boost::safe_numbers::ilog10(a); }); +} diff --git a/test/test_ilog2_sycl.cpp b/test/test_ilog2_sycl.cpp new file mode 100644 index 0000000..c0c6d16 --- /dev/null +++ b/test/test_ilog2_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("ilog2", [](auto a) { return boost::safe_numbers::ilog2(a); }); +} diff --git a/test/test_ilog_sycl.cpp b/test/test_ilog_sycl.cpp new file mode 100644 index 0000000..b724e25 --- /dev/null +++ b/test/test_ilog_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("ilog", [](auto a) { return boost::safe_numbers::ilog(a, decltype(a){7}); }); +} diff --git a/test/test_ipow_sycl.cpp b/test/test_ipow_sycl.cpp new file mode 100644 index 0000000..0b5f272 --- /dev/null +++ b/test/test_ipow_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_binary_all_unsigned("ipow", [](auto a, auto b) { return boost::safe_numbers::ipow(a, b); }); +} diff --git a/test/test_is_power_10_sycl.cpp b/test/test_is_power_10_sycl.cpp new file mode 100644 index 0000000..fc8aa62 --- /dev/null +++ b/test/test_is_power_10_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("is_power_10", [](auto a) { return static_cast(boost::safe_numbers::is_power_10(a)); }); +} diff --git a/test/test_is_power_2_sycl.cpp b/test/test_is_power_2_sycl.cpp new file mode 100644 index 0000000..7d2a15f --- /dev/null +++ b/test/test_is_power_2_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("is_power_2", [](auto a) { return static_cast(boost::safe_numbers::is_power_2(a)); }); +} diff --git a/test/test_isqrt_sycl.cpp b/test/test_isqrt_sycl.cpp new file mode 100644 index 0000000..0877f92 --- /dev/null +++ b/test/test_isqrt_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("isqrt", [](auto a) { return boost::safe_numbers::isqrt(a); }); +} diff --git a/test/test_lcm_sycl.cpp b/test/test_lcm_sycl.cpp new file mode 100644 index 0000000..25d5539 --- /dev/null +++ b/test/test_lcm_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_binary_all_unsigned("lcm", [](auto a, auto b) { return boost::safe_numbers::lcm(a, b); }); +} diff --git a/test/test_midpoint_sycl.cpp b/test/test_midpoint_sycl.cpp new file mode 100644 index 0000000..2049223 --- /dev/null +++ b/test/test_midpoint_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_binary_all_unsigned("midpoint", [](auto a, auto b) { return boost::safe_numbers::midpoint(a, b); }); +} diff --git a/test/test_next_multiple_of_sycl.cpp b/test/test_next_multiple_of_sycl.cpp new file mode 100644 index 0000000..e27639a --- /dev/null +++ b/test/test_next_multiple_of_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_binary_all_unsigned("next_multiple_of", [](auto a, auto b) { return boost::safe_numbers::next_multiple_of(a, b); }); +} diff --git a/test/test_popcount_sycl.cpp b/test/test_popcount_sycl.cpp new file mode 100644 index 0000000..703d697 --- /dev/null +++ b/test/test_popcount_sycl.cpp @@ -0,0 +1,11 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned( + "popcount", [](auto a) { return boost::safe_numbers::popcount(a); }); +} diff --git a/test/test_remove_trailing_zeros_sycl.cpp b/test/test_remove_trailing_zeros_sycl.cpp new file mode 100644 index 0000000..80cf3c0 --- /dev/null +++ b/test/test_remove_trailing_zeros_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("remove_trailing_zeros", [](auto a) { return decltype(a){boost::safe_numbers::remove_trailing_zeros(a).trimmed_number}; }); +} diff --git a/test/test_rotl_sycl.cpp b/test/test_rotl_sycl.cpp new file mode 100644 index 0000000..3fc89b8 --- /dev/null +++ b/test/test_rotl_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("rotl", [](auto a) { return boost::safe_numbers::rotl(a, 3); }); +} diff --git a/test/test_rotr_sycl.cpp b/test/test_rotr_sycl.cpp new file mode 100644 index 0000000..d2a5684 --- /dev/null +++ b/test/test_rotr_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("rotr", [](auto a) { return boost::safe_numbers::rotr(a, 3); }); +} diff --git a/test/test_signed_add_sycl.cpp b/test/test_signed_add_sycl.cpp new file mode 100644 index 0000000..aea2a4c --- /dev/null +++ b/test/test_signed_add_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_all_signed("add", [](auto a, auto b) { return a + b; }); +} diff --git a/test/test_signed_div_sycl.cpp b/test/test_signed_div_sycl.cpp new file mode 100644 index 0000000..a0a238b --- /dev/null +++ b/test/test_signed_div_sycl.cpp @@ -0,0 +1,11 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_all_signed("div", + [](auto a, auto b) { return a / sn_sycl_test::nonzero(b); }); +} diff --git a/test/test_signed_mod_sycl.cpp b/test/test_signed_mod_sycl.cpp new file mode 100644 index 0000000..f2a3ce7 --- /dev/null +++ b/test/test_signed_mod_sycl.cpp @@ -0,0 +1,11 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_all_signed("mod", + [](auto a, auto b) { return a % sn_sycl_test::nonzero(b); }); +} diff --git a/test/test_signed_mul_sycl.cpp b/test/test_signed_mul_sycl.cpp new file mode 100644 index 0000000..f67eb26 --- /dev/null +++ b/test/test_signed_mul_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_all_signed("mul", [](auto a, auto b) { return a * b; }); +} diff --git a/test/test_signed_ops_error_sycl.cpp b/test/test_signed_ops_error_sycl.cpp new file mode 100644 index 0000000..701fe1c --- /dev/null +++ b/test/test_signed_ops_error_sycl.cpp @@ -0,0 +1,62 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// Signed arithmetic error paths (no CUDA analog exists for these, but the +// signed family has error conditions the unsigned matrix cannot reach): +// add overflow/underflow, division and modulo by zero, the min / -1 +// overflow, and unary negation of min. + +#include "sycl_test.hpp" +#include + +template +int run(const char* name) +{ + using basis_type = typename T::basis_type; + constexpr auto max_val {(std::numeric_limits::max)()}; + constexpr auto min_val {(std::numeric_limits::min)()}; + + int rc {EXIT_SUCCESS}; + + rc |= sn_sycl_test::expect_binary_throw(name, + T{max_val}, T{max_val}, + [](auto x, auto y) { return x + y; }); + + rc |= sn_sycl_test::expect_binary_throw(name, + T{min_val}, T{min_val}, + [](auto x, auto y) { return x + y; }); + + rc |= sn_sycl_test::expect_binary_throw(name, + T{basis_type{1}}, T{basis_type{0}}, + [](auto x, auto y) { return x / y; }); + + rc |= sn_sycl_test::expect_binary_throw(name, + T{min_val}, T{static_cast(-1)}, + [](auto x, auto y) { return x / y; }); + + rc |= sn_sycl_test::expect_binary_throw(name, + T{basis_type{1}}, T{basis_type{0}}, + [](auto x, auto y) { return x % y; }); + + rc |= sn_sycl_test::expect_binary_throw(name, + T{min_val}, T{static_cast(-1)}, + [](auto x, auto y) { return x % y; }); + + rc |= sn_sycl_test::expect_binary_throw(name, + T{min_val}, T{basis_type{0}}, + [](auto x, auto y) { static_cast(y); return -x; }); + + return rc; +} + +int main() +{ + using namespace boost::safe_numbers; + + int rc {EXIT_SUCCESS}; + rc |= run("i8"); + rc |= run("i64"); + rc |= run("i128"); + return rc; +} diff --git a/test/test_signed_sub_sycl.cpp b/test/test_signed_sub_sycl.cpp new file mode 100644 index 0000000..28185ca --- /dev/null +++ b/test/test_signed_sub_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_all_signed("sub", [](auto a, auto b) { return a - b; }); +} diff --git a/test/test_to_be_sycl.cpp b/test/test_to_be_sycl.cpp new file mode 100644 index 0000000..5c6b83d --- /dev/null +++ b/test/test_to_be_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("to_be", [](auto a) { return boost::safe_numbers::to_be(a); }); +} diff --git a/test/test_to_le_sycl.cpp b/test/test_to_le_sycl.cpp new file mode 100644 index 0000000..5c5f843 --- /dev/null +++ b/test/test_to_le_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_unary_all_unsigned("to_le", [](auto a) { return boost::safe_numbers::to_le(a); }); +} diff --git a/test/test_unsigned_add_error_sycl.cpp b/test/test_unsigned_add_error_sycl.cpp new file mode 100644 index 0000000..053b01f --- /dev/null +++ b/test/test_unsigned_add_error_sycl.cpp @@ -0,0 +1,32 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// An unsigned addition that overflows must be reported from the device and surface on the host as +// std::overflow_error via device_error_context::synchronize(). + +#include "sycl_test.hpp" +#include + +template +int run(const char* label) +{ + using basis_type = typename T::basis_type; + + return sn_sycl_test::expect_binary_throw(label, + T{(std::numeric_limits::max)()}, T{(std::numeric_limits::max)()}, + [](auto x, auto y) { return x + y; }); +} + +int main() +{ + using namespace boost::safe_numbers; + + int rc {EXIT_SUCCESS}; + rc |= run("u8 add"); + rc |= run("u16 add"); + rc |= run("u32 add"); + rc |= run("u64 add"); + rc |= run("u128 add"); + return rc; +} diff --git a/test/test_unsigned_add_sycl.cpp b/test/test_unsigned_add_sycl.cpp new file mode 100644 index 0000000..f2a35a0 --- /dev/null +++ b/test/test_unsigned_add_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_all_unsigned("add", [](auto a, auto b) { return a + b; }); +} diff --git a/test/test_unsigned_div_error_sycl.cpp b/test/test_unsigned_div_error_sycl.cpp new file mode 100644 index 0000000..76b9711 --- /dev/null +++ b/test/test_unsigned_div_error_sycl.cpp @@ -0,0 +1,32 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// An unsigned division by zero must be reported from the device and surface on the host as +// std::domain_error via device_error_context::synchronize(). + +#include "sycl_test.hpp" +#include + +template +int run(const char* label) +{ + using basis_type = typename T::basis_type; + + return sn_sycl_test::expect_binary_throw(label, + T{basis_type{1}}, T{basis_type{0}}, + [](auto x, auto y) { return x / y; }); +} + +int main() +{ + using namespace boost::safe_numbers; + + int rc {EXIT_SUCCESS}; + rc |= run("u8 div"); + rc |= run("u16 div"); + rc |= run("u32 div"); + rc |= run("u64 div"); + rc |= run("u128 div"); + return rc; +} diff --git a/test/test_unsigned_div_sycl.cpp b/test/test_unsigned_div_sycl.cpp new file mode 100644 index 0000000..65f84ed --- /dev/null +++ b/test/test_unsigned_div_sycl.cpp @@ -0,0 +1,11 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_all_unsigned("div", + [](auto a, auto b) { return a / sn_sycl_test::nonzero(b); }); +} diff --git a/test/test_unsigned_mod_error_sycl.cpp b/test/test_unsigned_mod_error_sycl.cpp new file mode 100644 index 0000000..1085687 --- /dev/null +++ b/test/test_unsigned_mod_error_sycl.cpp @@ -0,0 +1,32 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// An unsigned modulo by zero must be reported from the device and surface on the host as +// std::domain_error via device_error_context::synchronize(). + +#include "sycl_test.hpp" +#include + +template +int run(const char* label) +{ + using basis_type = typename T::basis_type; + + return sn_sycl_test::expect_binary_throw(label, + T{basis_type{1}}, T{basis_type{0}}, + [](auto x, auto y) { return x % y; }); +} + +int main() +{ + using namespace boost::safe_numbers; + + int rc {EXIT_SUCCESS}; + rc |= run("u8 mod"); + rc |= run("u16 mod"); + rc |= run("u32 mod"); + rc |= run("u64 mod"); + rc |= run("u128 mod"); + return rc; +} diff --git a/test/test_unsigned_mod_sycl.cpp b/test/test_unsigned_mod_sycl.cpp new file mode 100644 index 0000000..474d53d --- /dev/null +++ b/test/test_unsigned_mod_sycl.cpp @@ -0,0 +1,11 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_all_unsigned("mod", + [](auto a, auto b) { return a % sn_sycl_test::nonzero(b); }); +} diff --git a/test/test_unsigned_mul_error_sycl.cpp b/test/test_unsigned_mul_error_sycl.cpp new file mode 100644 index 0000000..0560b31 --- /dev/null +++ b/test/test_unsigned_mul_error_sycl.cpp @@ -0,0 +1,32 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// An unsigned multiplication that overflows must be reported from the device and surface on the host as +// std::overflow_error via device_error_context::synchronize(). + +#include "sycl_test.hpp" +#include + +template +int run(const char* label) +{ + using basis_type = typename T::basis_type; + + return sn_sycl_test::expect_binary_throw(label, + T{(std::numeric_limits::max)()}, T{basis_type{2}}, + [](auto x, auto y) { return x * y; }); +} + +int main() +{ + using namespace boost::safe_numbers; + + int rc {EXIT_SUCCESS}; + rc |= run("u8 mul"); + rc |= run("u16 mul"); + rc |= run("u32 mul"); + rc |= run("u64 mul"); + rc |= run("u128 mul"); + return rc; +} diff --git a/test/test_unsigned_mul_sycl.cpp b/test/test_unsigned_mul_sycl.cpp new file mode 100644 index 0000000..34d42fa --- /dev/null +++ b/test/test_unsigned_mul_sycl.cpp @@ -0,0 +1,10 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + return sn_sycl_test::run_all_unsigned("mul", [](auto a, auto b) { return a * b; }); +} diff --git a/test/test_unsigned_sub_error_sycl.cpp b/test/test_unsigned_sub_error_sycl.cpp new file mode 100644 index 0000000..eb1873b --- /dev/null +++ b/test/test_unsigned_sub_error_sycl.cpp @@ -0,0 +1,32 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// An unsigned subtraction that underflows must be reported from the device and surface on the host as +// std::underflow_error via device_error_context::synchronize(). + +#include "sycl_test.hpp" +#include + +template +int run(const char* label) +{ + using basis_type = typename T::basis_type; + + return sn_sycl_test::expect_binary_throw(label, + T{basis_type{0}}, T{basis_type{1}}, + [](auto x, auto y) { return x - y; }); +} + +int main() +{ + using namespace boost::safe_numbers; + + int rc {EXIT_SUCCESS}; + rc |= run("u8 sub"); + rc |= run("u16 sub"); + rc |= run("u32 sub"); + rc |= run("u64 sub"); + rc |= run("u128 sub"); + return rc; +} diff --git a/test/test_unsigned_sub_sycl.cpp b/test/test_unsigned_sub_sycl.cpp new file mode 100644 index 0000000..b99f12b --- /dev/null +++ b/test/test_unsigned_sub_sycl.cpp @@ -0,0 +1,12 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt + +#include "sycl_test.hpp" + +int main() +{ + // Order operands so unsigned subtraction never underflows. + return sn_sycl_test::run_all_unsigned("sub", + [](auto a, auto b) { return a >= b ? a - b : b - a; }); +}