diff --git a/doc/modules/ROOT/nav.adoc b/doc/modules/ROOT/nav.adoc index b6aecab..a7a366a 100644 --- a/doc/modules/ROOT/nav.adoc +++ b/doc/modules/ROOT/nav.adoc @@ -23,6 +23,7 @@ ** xref:examples.adoc#examples_functional[Hashing] ** xref:examples.adoc#examples_random[Random Number Generation] ** xref:examples.adoc#examples_cuda[CUDA Device Support] +** xref:examples.adoc#examples_cuda_bounded[Bounded Types on a CUDA Device] ** 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] diff --git a/doc/modules/ROOT/pages/cuda.adoc b/doc/modules/ROOT/pages/cuda.adoc index c1487f7..edf3579 100644 --- a/doc/modules/ROOT/pages/cuda.adoc +++ b/doc/modules/ROOT/pages/cuda.adoc @@ -10,7 +10,7 @@ https://www.boost.org/LICENSE_1_0.txt == Description -All integer types of this library (unsigned `u8`-`u128`, signed `i8`-`i128`, and bounded types) support compilation with NVCC. +All integer types of this library (unsigned `u8`-`u128`, signed `i8`-`i128`, and bounded types) support compilation with NVCC and use inside device kernels. The safe floating-point types `f32` and `f64` are supported on the device as well, using the same error-reporting mechanism; see <>. To get the safety guarantees, there are some small modifications to the way that CUDA code is written. Normally you would have something like this: @@ -136,7 +136,9 @@ The mapping mirrors the host behavior of the types documented in xref:floats.ado |=== The `device_exception_mode` and the choice between `trapped` and `untrapped` apply unchanged, so the same `device_error_context` catches overflow, underflow, and domain errors from `f32` and `f64` kernels. -The range-checked `bounded_float` type (see xref:bounded_float.adoc[]) validates its value at construction and is a host-side type, so the device-enabled floating-point types are `f32` and `f64`. +The range-checked `bounded_float` type runs on the device as well; a result that violates its IEEE 754 contract or leaves the `[Min, Max]` range is captured by the same `device_error_context` and rethrown on the host. +The bounded integer types `bounded_uint` and `bounded_int` are likewise device-enabled. +See xref:examples.adoc#examples_cuda_bounded[the bounded-types CUDA example]. [#cuda_device_exception_mode] == The `device_exception_mode` Enum @@ -282,6 +284,9 @@ If an error was captured by device code, the error state is cleared and the appr | Domain error (e.g. division by zero) | `std::domain_error` +| Invalid argument +| `std::invalid_argument` + | Unknown | `std::runtime_error` |=== diff --git a/doc/modules/ROOT/pages/examples.adoc b/doc/modules/ROOT/pages/examples.adoc index c397e8f..8aeea6a 100644 --- a/doc/modules/ROOT/pages/examples.adoc +++ b/doc/modules/ROOT/pages/examples.adoc @@ -771,6 +771,28 @@ Charconv (rt): PASSED ---- ==== +[#examples_cuda_bounded] +== Bounded Types on a CUDA Device + +The range-checked bounded types (`bounded_uint`, `bounded_int`, and `bounded_float`) are also annotated with `\\__host__ __device__` and run on a CUDA device. +In-range arithmetic matches the host result, and a result that leaves the `[Min, Max]` range is captured by `device_error_context` and rethrown as `std::domain_error` on the host. + +.This https://github.com/boostorg/safe_numbers/blob/develop/examples/cuda_bounded.cu[example] demonstrates the bounded types running on a CUDA device. +==== +[source, c++] +---- +include::example$cuda_bounded.cu[] +---- + +Output: +---- +bounded_uint add: PASSED +bounded_int mul: PASSED +bounded_float add: PASSED +out-of-range sum: caught std::domain_error +---- +==== + [#examples_cuda_error_handling] == CUDA Error Handling diff --git a/examples/cuda_bounded.cu b/examples/cuda_bounded.cu new file mode 100644 index 0000000..18d136d --- /dev/null +++ b/examples/cuda_bounded.cu @@ -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 + +// This example demonstrates the range-checked bounded types +// (bounded_uint, bounded_int, bounded_float) running on a CUDA device. +// In-range arithmetic is verified against the host result, and an +// out-of-range result is caught on the host through device_error_context. + +#include +#include +#include +#include +#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>; + +__global__ void uint_add(const bu* a, const bu* b, bu* out) { *out = *a + *b; } +__global__ void int_mul(const bi* a, const bi* b, bi* out) { *out = *a * *b; } +__global__ void float_add(const bf* a, const bf* b, bf* out) { *out = *a + *b; } + +template +T* managed() +{ + T* p = nullptr; + cudaMallocManaged(&p, sizeof(T)); + cudaDeviceSynchronize(); + return p; +} + +int main() +{ + // untrapped mode: a device error is deferred to the host rather than + // trapping the kernel, so the context can be reused after catching. + boost::safe_numbers::device_error_context ctx {boost::safe_numbers::untrapped}; + + // bounded_uint: 300000 + 400000 = 700000, within [0, 1000000] + bu* ua {managed()}; + bu* ub {managed()}; + bu* uo {managed()}; + *ua = bu{300000}; + *ub = bu{400000}; + uint_add<<<1, 1>>>(ua, ub, uo); + ctx.synchronize(); + std::cout << "bounded_uint add: " << ((*uo == *ua + *ub) ? "PASSED" : "FAILED") << std::endl; + + // bounded_int: -600 * 700 = -420000, within [-1000000, 1000000] + bi* ia {managed()}; + bi* ib {managed()}; + bi* io {managed()}; + *ia = bi{-600}; + *ib = bi{700}; + int_mul<<<1, 1>>>(ia, ib, io); + ctx.synchronize(); + std::cout << "bounded_int mul: " << ((*io == *ia * *ib) ? "PASSED" : "FAILED") << std::endl; + + // bounded_float: 250.5 + 749.25 = 999.75, within [0, 1000000] + bf* fa {managed()}; + bf* fb {managed()}; + bf* fo {managed()}; + *fa = bf{bf::basis_type{250.5f}}; + *fb = bf{bf::basis_type{749.25f}}; + float_add<<<1, 1>>>(fa, fb, fo); + ctx.synchronize(); + std::cout << "bounded_float add: " << ((*fo == *fa + *fb) ? "PASSED" : "FAILED") << std::endl; + + // Out of range: 800000 + 800000 = 1600000 exceeds the [0, 1000000] bound + *ua = bu{800000}; + *ub = bu{800000}; + uint_add<<<1, 1>>>(ua, ub, uo); + try + { + ctx.synchronize(); + std::cout << "out-of-range: no error thrown (unexpected)" << std::endl; + } + catch (const std::domain_error&) + { + std::cout << "out-of-range sum: caught std::domain_error" << std::endl; + } + + cudaFree(ua); + cudaFree(ub); + cudaFree(uo); + cudaFree(ia); + cudaFree(ib); + cudaFree(io); + cudaFree(fa); + cudaFree(fb); + cudaFree(fo); + + return 0; +} diff --git a/include/boost/safe_numbers/bounded_floats.hpp b/include/boost/safe_numbers/bounded_floats.hpp index 3b07603..374235c 100644 --- a/include/boost/safe_numbers/bounded_floats.hpp +++ b/include/boost/safe_numbers/bounded_floats.hpp @@ -54,18 +54,20 @@ class bounded_float public: - explicit constexpr bounded_float(const basis_type val) + BOOST_SAFE_NUMBERS_HOST_DEVICE explicit constexpr bounded_float(const basis_type val) { const auto raw {static_cast(val)}; // 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 (std::is_constant_evaluated()) { throw std::domain_error("bounded_float NaN value"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_float NaN value"); } @@ -76,11 +78,13 @@ class bounded_float if (raw < min_raw || raw > max_raw) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_float value out of range"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_float value out of range"); } @@ -89,11 +93,11 @@ class bounded_float basis_ = val; } - explicit constexpr bounded_float(const underlying_type val) : bounded_float{basis_type{val}} {} + BOOST_SAFE_NUMBERS_HOST_DEVICE explicit constexpr bounded_float(const underlying_type val) : bounded_float{basis_type{val}} {} template requires (detail::is_compatible_float_type) - [[nodiscard]] explicit constexpr operator OtherBasis() const + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] explicit constexpr operator OtherBasis() const { const auto raw {static_cast(basis_)}; @@ -103,11 +107,13 @@ class bounded_float if (detail::impl::constexpr_isinf(result) && !detail::impl::constexpr_isinf(raw)) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::overflow_error("bounded_float narrowing conversion overflow"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::overflow_error, "bounded_float narrowing conversion overflow"); } @@ -122,7 +128,7 @@ class bounded_float } template - [[nodiscard]] explicit constexpr operator bounded_float() const + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] explicit constexpr operator bounded_float() const { using target_basis = typename bounded_float::basis_type; using target_underlying = detail::underlying_type_t; @@ -130,20 +136,20 @@ class bounded_float return bounded_float{target_basis{static_cast(raw)}}; } - [[nodiscard]] constexpr auto to_basis() const noexcept -> basis_type { return basis_; } + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto to_basis() const noexcept -> basis_type { return basis_; } - [[nodiscard]] friend constexpr auto operator==(bounded_float lhs, bounded_float rhs) noexcept -> bool = default; + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] friend constexpr auto operator==(bounded_float lhs, bounded_float rhs) noexcept -> bool = default; - [[nodiscard]] friend constexpr auto operator<=>(bounded_float lhs, bounded_float rhs) noexcept + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] friend constexpr auto operator<=>(bounded_float lhs, bounded_float rhs) noexcept -> std::partial_ordering = default; - constexpr auto operator+=(bounded_float rhs) -> bounded_float&; + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator+=(bounded_float rhs) -> bounded_float&; - constexpr auto operator-=(bounded_float rhs) -> bounded_float&; + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator-=(bounded_float rhs) -> bounded_float&; - constexpr auto operator*=(bounded_float rhs) -> bounded_float&; + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator*=(bounded_float rhs) -> bounded_float&; - constexpr auto operator/=(bounded_float rhs) -> bounded_float&; + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator/=(bounded_float rhs) -> bounded_float&; }; // ------------------------------ @@ -155,7 +161,7 @@ class bounded_float // The bounded_float constructor then re-validates the result against [Min, Max]. template -[[nodiscard]] constexpr auto operator+(const bounded_float lhs, +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator+(const bounded_float lhs, const bounded_float rhs) -> bounded_float { using basis = typename bounded_float::basis_type; @@ -166,7 +172,7 @@ template } template -[[nodiscard]] constexpr auto operator-(const bounded_float lhs, +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator-(const bounded_float lhs, const bounded_float rhs) -> bounded_float { using basis = typename bounded_float::basis_type; @@ -177,7 +183,7 @@ template } template -[[nodiscard]] constexpr auto operator*(const bounded_float lhs, +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator*(const bounded_float lhs, const bounded_float rhs) -> bounded_float { using basis = typename bounded_float::basis_type; @@ -188,7 +194,7 @@ template } template -[[nodiscard]] constexpr auto operator/(const bounded_float lhs, +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator/(const bounded_float lhs, const bounded_float rhs) -> bounded_float { using basis = typename bounded_float::basis_type; @@ -209,7 +215,7 @@ template detail::float_raw_value(Min) == detail::float_raw_value(Min) && detail::float_raw_value(Max) == detail::float_raw_value(Max) && detail::float_raw_value(Max) > detail::float_raw_value(Min)) -constexpr auto bounded_float::operator+=(bounded_float rhs) -> bounded_float& +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_float::operator+=(bounded_float rhs) -> bounded_float& { *this = *this + rhs; return *this; @@ -222,7 +228,7 @@ template detail::float_raw_value(Min) == detail::float_raw_value(Min) && detail::float_raw_value(Max) == detail::float_raw_value(Max) && detail::float_raw_value(Max) > detail::float_raw_value(Min)) -constexpr auto bounded_float::operator-=(bounded_float rhs) -> bounded_float& +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_float::operator-=(bounded_float rhs) -> bounded_float& { *this = *this - rhs; return *this; @@ -235,7 +241,7 @@ template detail::float_raw_value(Min) == detail::float_raw_value(Min) && detail::float_raw_value(Max) == detail::float_raw_value(Max) && detail::float_raw_value(Max) > detail::float_raw_value(Min)) -constexpr auto bounded_float::operator*=(bounded_float rhs) -> bounded_float& +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_float::operator*=(bounded_float rhs) -> bounded_float& { *this = *this * rhs; return *this; @@ -248,7 +254,7 @@ template detail::float_raw_value(Min) == detail::float_raw_value(Min) && detail::float_raw_value(Max) == detail::float_raw_value(Max) && detail::float_raw_value(Max) > detail::float_raw_value(Min)) -constexpr auto bounded_float::operator/=(bounded_float rhs) -> bounded_float& +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_float::operator/=(bounded_float rhs) -> bounded_float& { *this = *this / rhs; return *this; diff --git a/include/boost/safe_numbers/bounded_integers.hpp b/include/boost/safe_numbers/bounded_integers.hpp index c970e71..4b0db49 100644 --- a/include/boost/safe_numbers/bounded_integers.hpp +++ b/include/boost/safe_numbers/bounded_integers.hpp @@ -48,6 +48,22 @@ auto to_string_val(T val) -> std::string return to_string(val); } +// Range checks that take the bounds as runtime parameters. A bound of 0 reaches +// the comparison as a parameter rather than a literal, so an unsigned check does +// not trip the "pointless comparison of unsigned integer with zero" diagnostic +// (for example NVCC warning 186) when Min is 0. +template +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto out_of_bounds(const U value, const U min_val, const U max_val) noexcept -> bool +{ + return value < min_val || value > max_val; +} + +template +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto below_bound(const U value, const U min_val) noexcept -> bool +{ + return value < min_val; +} + } // namespace detail template @@ -70,7 +86,7 @@ class bounded_uint public: - explicit constexpr bounded_uint(const basis_type val) + BOOST_SAFE_NUMBERS_HOST_DEVICE explicit constexpr bounded_uint(const basis_type val) { constexpr auto min_raw {static_cast(detail::raw_value(Min))}; constexpr auto max_raw {static_cast(detail::raw_value(Max))}; @@ -79,11 +95,13 @@ class bounded_uint if (val < min_val || val > max_val) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint value out of range"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_uint value out of range"); } @@ -92,11 +110,11 @@ class bounded_uint basis_ = val; } - explicit constexpr bounded_uint(const underlying_type val) : bounded_uint{basis_type{val}} {} + BOOST_SAFE_NUMBERS_HOST_DEVICE explicit constexpr bounded_uint(const underlying_type val) : bounded_uint{basis_type{val}} {} template requires (detail::is_unsigned_library_type_v || detail::is_fundamental_unsigned_integral_v) - [[nodiscard]] explicit constexpr operator OtherBasis() const + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] explicit constexpr operator OtherBasis() const { const auto raw {static_cast>(basis_)}; @@ -105,11 +123,13 @@ 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 (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint conversion overflow"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_uint conversion overflow"); } @@ -124,40 +144,40 @@ class bounded_uint } template - [[nodiscard]] explicit constexpr operator bounded_uint() const + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] explicit constexpr operator bounded_uint() const { using target_basis = typename bounded_uint::basis_type; const auto raw {static_cast>(basis_)}; return bounded_uint{static_cast(raw)}; } - [[nodiscard]] explicit constexpr operator basis_type() const noexcept { return basis_; } + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] explicit constexpr operator basis_type() const noexcept { return basis_; } - [[nodiscard]] explicit constexpr operator underlying_type() const noexcept { return static_cast(basis_); } + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] explicit constexpr operator underlying_type() const noexcept { return static_cast(basis_); } - [[nodiscard]] friend constexpr auto operator<=>(bounded_uint lhs, bounded_uint rhs) noexcept + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] friend constexpr auto operator<=>(bounded_uint lhs, bounded_uint rhs) noexcept -> std::strong_ordering = default; - constexpr auto operator+=(bounded_uint rhs) -> bounded_uint&; + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator+=(bounded_uint rhs) -> bounded_uint&; - constexpr auto operator-=(bounded_uint rhs) -> bounded_uint&; + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator-=(bounded_uint rhs) -> bounded_uint&; - constexpr auto operator*=(bounded_uint rhs) -> bounded_uint&; + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator*=(bounded_uint rhs) -> bounded_uint&; - constexpr auto operator/=(bounded_uint rhs) -> bounded_uint&; + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator/=(bounded_uint rhs) -> bounded_uint&; - constexpr auto operator++() -> bounded_uint&; + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator++() -> bounded_uint&; - constexpr auto operator++(int) -> bounded_uint; + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator++(int) -> bounded_uint; - constexpr auto operator--() -> bounded_uint&; + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator--() -> bounded_uint&; - constexpr auto operator--(int) -> bounded_uint; + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator--(int) -> bounded_uint; }; template -[[nodiscard]] constexpr auto operator+(const bounded_uint lhs, +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator+(const bounded_uint lhs, const bounded_uint rhs) -> bounded_uint { using basis = typename bounded_uint::basis_type; @@ -170,23 +190,27 @@ template underlying res {}; if (detail::impl::unsigned_no_intrin_add(lhs_raw, rhs_raw, res)) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::overflow_error("bounded_uint addition overflow"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::overflow_error, "bounded_uint addition overflow"); } } - if (res < min_raw || res > max_raw) + if (detail::out_of_bounds(res, min_raw, max_raw)) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint addition result out of range"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_uint addition result out of range"); } @@ -196,7 +220,7 @@ template } template -[[nodiscard]] constexpr auto operator-(const bounded_uint lhs, +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator-(const bounded_uint lhs, const bounded_uint rhs) -> bounded_uint { using basis = typename bounded_uint::basis_type; @@ -209,23 +233,27 @@ template underlying res {}; if (detail::impl::unsigned_no_intrin_sub(lhs_raw, rhs_raw, res)) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::underflow_error("bounded_uint subtraction underflow"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::underflow_error, "bounded_uint subtraction underflow"); } } - if (res < min_raw || res > max_raw) + if (detail::out_of_bounds(res, min_raw, max_raw)) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint subtraction result out of range"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_uint subtraction result out of range"); } @@ -235,7 +263,7 @@ template } template -[[nodiscard]] constexpr auto operator*(const bounded_uint lhs, +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator*(const bounded_uint lhs, const bounded_uint rhs) -> bounded_uint { using basis = typename bounded_uint::basis_type; @@ -248,23 +276,27 @@ template underlying res {}; if (detail::impl::no_intrin_mul(lhs_raw, rhs_raw, res)) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::overflow_error("bounded_uint multiplication overflow"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::overflow_error, "bounded_uint multiplication overflow"); } } - if (res < min_raw || res > max_raw) + if (detail::out_of_bounds(res, min_raw, max_raw)) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint multiplication result out of range"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_uint multiplication result out of range"); } @@ -274,7 +306,7 @@ template } template -[[nodiscard]] constexpr auto operator/(const bounded_uint lhs, +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator/(const bounded_uint lhs, const bounded_uint rhs) -> bounded_uint { using basis = typename bounded_uint::basis_type; @@ -286,11 +318,13 @@ template if (rhs_raw == 0U) [[unlikely]] { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint division by zero"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_uint division by zero"); } @@ -306,13 +340,15 @@ template res = lhs_raw / rhs_raw; } - if (res < min_raw || res > max_raw) + if (detail::out_of_bounds(res, min_raw, max_raw)) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint division result out of range"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_uint division result out of range"); } @@ -322,7 +358,7 @@ template } template -[[nodiscard]] constexpr auto operator%(const bounded_uint lhs, +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator%(const bounded_uint lhs, const bounded_uint rhs) -> bounded_uint { using basis = typename bounded_uint::basis_type; @@ -334,11 +370,13 @@ template if (rhs_raw == 0U) [[unlikely]] { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint modulo by zero"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_uint modulo by zero"); } @@ -354,13 +392,15 @@ template res = lhs_raw % rhs_raw; } - if (res < min_raw || res > max_raw) + if (detail::out_of_bounds(res, min_raw, max_raw)) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint modulo result out of range"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_uint modulo result out of range"); } @@ -373,7 +413,7 @@ template requires (detail::valid_bound && detail::valid_bound && detail::raw_value(Max) > detail::raw_value(Min)) -constexpr auto bounded_uint::operator+=(bounded_uint rhs) -> bounded_uint& +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_uint::operator+=(bounded_uint rhs) -> bounded_uint& { *this = *this + rhs; return *this; @@ -383,7 +423,7 @@ template requires (detail::valid_bound && detail::valid_bound && detail::raw_value(Max) > detail::raw_value(Min)) -constexpr auto bounded_uint::operator-=(bounded_uint rhs) -> bounded_uint& +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_uint::operator-=(bounded_uint rhs) -> bounded_uint& { *this = *this - rhs; return *this; @@ -393,7 +433,7 @@ template requires (detail::valid_bound && detail::valid_bound && detail::raw_value(Max) > detail::raw_value(Min)) -constexpr auto bounded_uint::operator*=(bounded_uint rhs) -> bounded_uint& +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_uint::operator*=(bounded_uint rhs) -> bounded_uint& { *this = *this * rhs; return *this; @@ -403,7 +443,7 @@ template requires (detail::valid_bound && detail::valid_bound && detail::raw_value(Max) > detail::raw_value(Min)) -constexpr auto bounded_uint::operator/=(bounded_uint rhs) -> bounded_uint& +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_uint::operator/=(bounded_uint rhs) -> bounded_uint& { *this = *this / rhs; return *this; @@ -413,7 +453,7 @@ template requires (detail::valid_bound && detail::valid_bound && detail::raw_value(Max) > detail::raw_value(Min)) -constexpr auto bounded_uint::operator++() -> bounded_uint& +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_uint::operator++() -> bounded_uint& { using underlying = detail::underlying_type_t; constexpr auto max_raw {static_cast(detail::raw_value(Max))}; @@ -422,11 +462,13 @@ constexpr auto bounded_uint::operator++() -> bounded_uint& underlying res {}; if (detail::impl::unsigned_no_intrin_add(raw, static_cast(1U), res)) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::overflow_error("bounded_uint increment overflow"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::overflow_error, "bounded_uint increment overflow"); } @@ -434,11 +476,13 @@ constexpr auto bounded_uint::operator++() -> bounded_uint& if (res > max_raw) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint increment result out of range"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_uint increment result out of range"); } @@ -452,7 +496,7 @@ template requires (detail::valid_bound && detail::valid_bound && detail::raw_value(Max) > detail::raw_value(Min)) -constexpr auto bounded_uint::operator++(int) -> bounded_uint +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_uint::operator++(int) -> bounded_uint { auto tmp {*this}; ++(*this); @@ -463,7 +507,7 @@ template requires (detail::valid_bound && detail::valid_bound && detail::raw_value(Max) > detail::raw_value(Min)) -constexpr auto bounded_uint::operator--() -> bounded_uint& +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_uint::operator--() -> bounded_uint& { using underlying = detail::underlying_type_t; constexpr auto min_raw {static_cast(detail::raw_value(Min))}; @@ -472,23 +516,27 @@ constexpr auto bounded_uint::operator--() -> bounded_uint& underlying res {}; if (detail::impl::unsigned_no_intrin_sub(raw, static_cast(1U), res)) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::underflow_error("bounded_uint decrement underflow"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::underflow_error, "bounded_uint decrement underflow"); } } - if (res < min_raw) + if (detail::below_bound(res, min_raw)) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_uint decrement result out of range"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_uint decrement result out of range"); } @@ -502,7 +550,7 @@ template requires (detail::valid_bound && detail::valid_bound && detail::raw_value(Max) > detail::raw_value(Min)) -constexpr auto bounded_uint::operator--(int) -> bounded_uint +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_uint::operator--(int) -> bounded_uint { auto tmp {*this}; --(*this); @@ -545,7 +593,7 @@ class bounded_int public: - explicit constexpr bounded_int(const basis_type val) + BOOST_SAFE_NUMBERS_HOST_DEVICE explicit constexpr bounded_int(const basis_type val) { constexpr auto min_raw {static_cast(detail::signed_raw_value(Min))}; constexpr auto max_raw {static_cast(detail::signed_raw_value(Max))}; @@ -554,11 +602,13 @@ class bounded_int if (val < min_val || val > max_val) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int value out of range"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_int value out of range"); } @@ -567,11 +617,11 @@ class bounded_int basis_ = val; } - explicit constexpr bounded_int(const underlying_type val) : bounded_int{basis_type{val}} {} + BOOST_SAFE_NUMBERS_HOST_DEVICE explicit constexpr bounded_int(const underlying_type val) : bounded_int{basis_type{val}} {} template requires (detail::is_signed_library_type_v || detail::is_fundamental_signed_integral_v) - [[nodiscard]] explicit constexpr operator OtherBasis() const + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] explicit constexpr operator OtherBasis() const { const auto raw {static_cast>(basis_)}; @@ -581,11 +631,13 @@ 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 (std::is_constant_evaluated()) { throw std::domain_error("bounded_int conversion overflow"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_int conversion overflow"); } @@ -600,7 +652,7 @@ class bounded_int } template - [[nodiscard]] explicit constexpr operator bounded_int() const + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] explicit constexpr operator bounded_int() const { using target_basis = typename bounded_int::basis_type; using target_underlying = detail::underlying_type_t; @@ -608,32 +660,32 @@ class bounded_int return bounded_int{target_basis{static_cast(raw)}}; } - [[nodiscard]] explicit constexpr operator basis_type() const noexcept { return basis_; } + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] explicit constexpr operator basis_type() const noexcept { return basis_; } - [[nodiscard]] explicit constexpr operator underlying_type() const noexcept { return static_cast(basis_); } + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] explicit constexpr operator underlying_type() const noexcept { return static_cast(basis_); } - [[nodiscard]] friend constexpr auto operator<=>(bounded_int lhs, bounded_int rhs) noexcept + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] friend constexpr auto operator<=>(bounded_int lhs, bounded_int rhs) noexcept -> std::strong_ordering = default; - [[nodiscard]] constexpr auto operator+() const noexcept -> bounded_int { return *this; } + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator+() const noexcept -> bounded_int { return *this; } - [[nodiscard]] constexpr auto operator-() const -> bounded_int; + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator-() const -> bounded_int; - constexpr auto operator+=(bounded_int rhs) -> bounded_int&; + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator+=(bounded_int rhs) -> bounded_int&; - constexpr auto operator-=(bounded_int rhs) -> bounded_int&; + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator-=(bounded_int rhs) -> bounded_int&; - constexpr auto operator*=(bounded_int rhs) -> bounded_int&; + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator*=(bounded_int rhs) -> bounded_int&; - constexpr auto operator/=(bounded_int rhs) -> bounded_int&; + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator/=(bounded_int rhs) -> bounded_int&; - constexpr auto operator++() -> bounded_int&; + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator++() -> bounded_int&; - constexpr auto operator++(int) -> bounded_int; + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator++(int) -> bounded_int; - constexpr auto operator--() -> bounded_int&; + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator--() -> bounded_int&; - constexpr auto operator--(int) -> bounded_int; + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator--(int) -> bounded_int; }; // ------------------------------ @@ -644,18 +696,20 @@ template requires (detail::valid_signed_bound && detail::valid_signed_bound && detail::signed_raw_value(Max) > detail::signed_raw_value(Min)) -constexpr auto bounded_int::operator-() const -> bounded_int +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_int::operator-() const -> bounded_int { using underlying = detail::underlying_type_t; const auto raw {static_cast(basis_)}; if (raw == std::numeric_limits::min()) [[unlikely]] { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::overflow_error("bounded_int negation overflow"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::overflow_error, "bounded_int negation overflow"); } @@ -670,7 +724,7 @@ constexpr auto bounded_int::operator-() const -> bounded_int // ------------------------------ template -[[nodiscard]] constexpr auto operator+(const bounded_int lhs, +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator+(const bounded_int lhs, const bounded_int rhs) -> bounded_int { using basis = typename bounded_int::basis_type; @@ -684,34 +738,40 @@ template 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 (std::is_constant_evaluated()) { throw std::overflow_error("bounded_int addition overflow"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::overflow_error, "bounded_int addition overflow"); } } else if (status == detail::impl::signed_overflow_status::underflow) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::underflow_error("bounded_int addition underflow"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::underflow_error, "bounded_int addition underflow"); } } - if (res < min_raw || res > max_raw) + if (detail::out_of_bounds(res, min_raw, max_raw)) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int addition result out of range"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_int addition result out of range"); } @@ -725,7 +785,7 @@ template // ------------------------------ template -[[nodiscard]] constexpr auto operator-(const bounded_int lhs, +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator-(const bounded_int lhs, const bounded_int rhs) -> bounded_int { using basis = typename bounded_int::basis_type; @@ -739,34 +799,40 @@ template 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 (std::is_constant_evaluated()) { throw std::overflow_error("bounded_int subtraction overflow"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::overflow_error, "bounded_int subtraction overflow"); } } else if (status == detail::impl::signed_overflow_status::underflow) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::underflow_error("bounded_int subtraction underflow"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::underflow_error, "bounded_int subtraction underflow"); } } - if (res < min_raw || res > max_raw) + if (detail::out_of_bounds(res, min_raw, max_raw)) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int subtraction result out of range"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_int subtraction result out of range"); } @@ -780,7 +846,7 @@ template // ------------------------------ template -[[nodiscard]] constexpr auto operator*(const bounded_int lhs, +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator*(const bounded_int lhs, const bounded_int rhs) -> bounded_int { using basis = typename bounded_int::basis_type; @@ -794,34 +860,40 @@ template 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 (std::is_constant_evaluated()) { throw std::overflow_error("bounded_int multiplication overflow"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::overflow_error, "bounded_int multiplication overflow"); } } else if (status == detail::impl::signed_overflow_status::underflow) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::underflow_error("bounded_int multiplication underflow"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::underflow_error, "bounded_int multiplication underflow"); } } - if (res < min_raw || res > max_raw) + if (detail::out_of_bounds(res, min_raw, max_raw)) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int multiplication result out of range"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_int multiplication result out of range"); } @@ -835,7 +907,7 @@ template // ------------------------------ template -[[nodiscard]] constexpr auto operator/(const bounded_int lhs, +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator/(const bounded_int lhs, const bounded_int rhs) -> bounded_int { using basis = typename bounded_int::basis_type; @@ -847,11 +919,13 @@ template if (rhs_raw == static_cast(0)) [[unlikely]] { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int division by zero"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_int division by zero"); } @@ -860,11 +934,13 @@ template if (lhs_raw == std::numeric_limits::min() && rhs_raw == static_cast(-1)) [[unlikely]] { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::overflow_error("bounded_int division overflow"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::overflow_error, "bounded_int division overflow"); } @@ -886,13 +962,15 @@ template res = lhs_raw / rhs_raw; } - if (res < min_raw || res > max_raw) + if (detail::out_of_bounds(res, min_raw, max_raw)) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int division result out of range"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_int division result out of range"); } @@ -906,7 +984,7 @@ template // ------------------------------ template -[[nodiscard]] constexpr auto operator%(const bounded_int lhs, +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator%(const bounded_int lhs, const bounded_int rhs) -> bounded_int { using basis = typename bounded_int::basis_type; @@ -918,11 +996,13 @@ template if (rhs_raw == static_cast(0)) [[unlikely]] { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int modulo by zero"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_int modulo by zero"); } @@ -931,11 +1011,13 @@ template if (lhs_raw == std::numeric_limits::min() && rhs_raw == static_cast(-1)) [[unlikely]] { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::overflow_error("bounded_int modulo overflow"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::overflow_error, "bounded_int modulo overflow"); } @@ -957,13 +1039,15 @@ template res = lhs_raw % rhs_raw; } - if (res < min_raw || res > max_raw) + if (detail::out_of_bounds(res, min_raw, max_raw)) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int modulo result out of range"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_int modulo result out of range"); } @@ -980,7 +1064,7 @@ template requires (detail::valid_signed_bound && detail::valid_signed_bound && detail::signed_raw_value(Max) > detail::signed_raw_value(Min)) -constexpr auto bounded_int::operator+=(bounded_int rhs) -> bounded_int& +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_int::operator+=(bounded_int rhs) -> bounded_int& { *this = *this + rhs; return *this; @@ -990,7 +1074,7 @@ template requires (detail::valid_signed_bound && detail::valid_signed_bound && detail::signed_raw_value(Max) > detail::signed_raw_value(Min)) -constexpr auto bounded_int::operator-=(bounded_int rhs) -> bounded_int& +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_int::operator-=(bounded_int rhs) -> bounded_int& { *this = *this - rhs; return *this; @@ -1000,7 +1084,7 @@ template requires (detail::valid_signed_bound && detail::valid_signed_bound && detail::signed_raw_value(Max) > detail::signed_raw_value(Min)) -constexpr auto bounded_int::operator*=(bounded_int rhs) -> bounded_int& +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_int::operator*=(bounded_int rhs) -> bounded_int& { *this = *this * rhs; return *this; @@ -1010,7 +1094,7 @@ template requires (detail::valid_signed_bound && detail::valid_signed_bound && detail::signed_raw_value(Max) > detail::signed_raw_value(Min)) -constexpr auto bounded_int::operator/=(bounded_int rhs) -> bounded_int& +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_int::operator/=(bounded_int rhs) -> bounded_int& { *this = *this / rhs; return *this; @@ -1024,7 +1108,7 @@ template requires (detail::valid_signed_bound && detail::valid_signed_bound && detail::signed_raw_value(Max) > detail::signed_raw_value(Min)) -constexpr auto bounded_int::operator++() -> bounded_int& +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_int::operator++() -> bounded_int& { using underlying = detail::underlying_type_t; constexpr auto max_raw {static_cast(detail::signed_raw_value(Max))}; @@ -1034,11 +1118,13 @@ constexpr auto bounded_int::operator++() -> bounded_int& 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 (std::is_constant_evaluated()) { throw std::overflow_error("bounded_int increment overflow"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::overflow_error, "bounded_int increment overflow"); } @@ -1046,11 +1132,13 @@ constexpr auto bounded_int::operator++() -> bounded_int& if (res > max_raw) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int increment result out of range"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_int increment result out of range"); } @@ -1064,7 +1152,7 @@ template requires (detail::valid_signed_bound && detail::valid_signed_bound && detail::signed_raw_value(Max) > detail::signed_raw_value(Min)) -constexpr auto bounded_int::operator++(int) -> bounded_int +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_int::operator++(int) -> bounded_int { auto tmp {*this}; ++(*this); @@ -1075,7 +1163,7 @@ template requires (detail::valid_signed_bound && detail::valid_signed_bound && detail::signed_raw_value(Max) > detail::signed_raw_value(Min)) -constexpr auto bounded_int::operator--() -> bounded_int& +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_int::operator--() -> bounded_int& { using underlying = detail::underlying_type_t; constexpr auto min_raw {static_cast(detail::signed_raw_value(Min))}; @@ -1085,23 +1173,27 @@ constexpr auto bounded_int::operator--() -> bounded_int& 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 (std::is_constant_evaluated()) { throw std::underflow_error("bounded_int decrement underflow"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::underflow_error, "bounded_int decrement underflow"); } } - if (res < min_raw) + if (detail::below_bound(res, min_raw)) { + #if !(defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) if (std::is_constant_evaluated()) { throw std::domain_error("bounded_int decrement result out of range"); // LCOV_EXCL_LINE } else + #endif { BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::domain_error, "bounded_int decrement result out of range"); } @@ -1115,7 +1207,7 @@ template requires (detail::valid_signed_bound && detail::valid_signed_bound && detail::signed_raw_value(Max) > detail::signed_raw_value(Min)) -constexpr auto bounded_int::operator--(int) -> bounded_int +BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_int::operator--(int) -> bounded_int { auto tmp {*this}; --(*this); diff --git a/include/boost/safe_numbers/cuda_error_reporting.hpp b/include/boost/safe_numbers/cuda_error_reporting.hpp index a0ba083..eaf8599 100644 --- a/include/boost/safe_numbers/cuda_error_reporting.hpp +++ b/include/boost/safe_numbers/cuda_error_reporting.hpp @@ -45,6 +45,7 @@ enum class exception_type : unsigned domain_error, overflow, underflow, + invalid_argument, unknown, }; @@ -76,7 +77,7 @@ BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto to_exception_enum() noexcept -> ex } else if constexpr (std::is_same_v) { - return exception_type::domain_error; + return exception_type::invalid_argument; } else { @@ -168,6 +169,9 @@ __host__ __device__ inline void report_device_error( 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: @@ -277,6 +281,9 @@ class device_error_context 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: diff --git a/test/Jamfile b/test/Jamfile index 6181eea..b228dd1 100644 --- a/test/Jamfile +++ b/test/Jamfile @@ -105,6 +105,7 @@ compile-fail compile_fail_unsigned_literals.cpp ; run test_signed_literals.cpp ; compile-fail compile_fail_signed_literals.cpp ; run test_literals.cpp ; +run test_device_error_mapping.cpp ; run test_unsigned_subtraction.cpp ; run test_unsigned_multiplication.cpp ; run test_unsigned_division.cpp ; diff --git a/test/cuda_jamfile b/test/cuda_jamfile index 5b6c91b..8990941 100644 --- a/test/cuda_jamfile +++ b/test/cuda_jamfile @@ -11,6 +11,7 @@ project : requirements # Test the error handler mechanisms run test_cuda_error_handling.cu ; +run test_cuda_invalid_argument.cu ; # u8 tests run test_cuda_u8_add.cu ; @@ -368,7 +369,34 @@ run test_cuda_f64_sub_error.cu ; run test_cuda_f64_mul_error.cu ; run test_cuda_f64_div_error.cu ; +# Bounded type tests + +# bounded_uint tests +run test_cuda_bounded_uint_add.cu ; +run test_cuda_bounded_uint_sub.cu ; +run test_cuda_bounded_uint_mul.cu ; +run test_cuda_bounded_uint_div.cu ; +run test_cuda_bounded_uint_add_error.cu ; + +# bounded_int tests +run test_cuda_bounded_int_add.cu ; +run test_cuda_bounded_int_sub.cu ; +run test_cuda_bounded_int_mul.cu ; +run test_cuda_bounded_int_div.cu ; +run test_cuda_bounded_int_add_error.cu ; + +# bounded_float tests +run test_cuda_bounded_float_add.cu ; +run test_cuda_bounded_float_sub.cu ; +run test_cuda_bounded_float_mul.cu ; +run test_cuda_bounded_float_div.cu ; +run test_cuda_bounded_float_add_error.cu ; + +# bounded mixed operators (modulo, increment/decrement, unary, compound assignment) +run test_cuda_bounded_ops.cu ; + # Examples run ../examples/cuda.cu ; run ../examples/cuda_error_handling.cu ; run ../examples/cuda_error_handling_without_error_context.cu ; +run ../examples/cuda_bounded.cu ; diff --git a/test/test_cuda_bounded_float_add.cu b/test/test_cuda_bounded_float_add.cu new file mode 100644 index 0000000..12c8830 --- /dev/null +++ b/test/test_cuda_bounded_float_add.cu @@ -0,0 +1,61 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" + +#include + +using test_type = boost::safe_numbers::bounded_float<0.0f, 1000000.0f>; +using basis_type = test_type::basis_type; + +__global__ void cuda_test(const test_type *a, const test_type *b, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + out[i] = a[i] + b[i]; + } +} + +int main() +{ + int numElements = 2048; + + cuda_managed_ptr A(numElements); + cuda_managed_ptr B(numElements); + cuda_managed_ptr OUT(numElements); + + std::mt19937_64 rng{42}; + std::uniform_real_distribution dist{0.0f, 400000.0f}; + for (int i = 0; i < numElements; ++i) + { + A[i] = test_type{basis_type{dist(rng)}}; + B[i] = test_type{basis_type{dist(rng)}}; + } + + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + + boost::safe_numbers::device_error_context ctx; + cuda_test<<>>(A.get(), B.get(), OUT.get(), numElements); + ctx.synchronize(); + + for (int i = 0; i < numElements; ++i) + { + if (OUT[i] != A[i] + B[i]) + { + std::cerr << "Result verification failed at element " << i << "!" << std::endl; + return EXIT_FAILURE; + } + } + + std::cout << "Test PASSED\n"; + return 0; +} diff --git a/test/test_cuda_bounded_float_add_error.cu b/test/test_cuda_bounded_float_add_error.cu new file mode 100644 index 0000000..b87050d --- /dev/null +++ b/test/test_cuda_bounded_float_add_error.cu @@ -0,0 +1,60 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" + +#include + +using test_type = boost::safe_numbers::bounded_float<0.0f, 1000.0f>; +using basis_type = test_type::basis_type; + +__global__ void cuda_test(const test_type *a, const test_type *b, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + out[i] = a[i] + b[i]; + } +} + +int main() +{ + int numElements = 1024; + + cuda_managed_ptr A(numElements); + cuda_managed_ptr B(numElements); + cuda_managed_ptr OUT(numElements); + + // 600 + 600 = 1200 exceeds the [0, 1000] bound -> out of range + for (int i = 0; i < numElements; ++i) + { + A[i] = test_type{basis_type{600.0f}}; + B[i] = test_type{basis_type{600.0f}}; + } + + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + + boost::safe_numbers::device_error_context ctx; + cuda_test<<>>(A.get(), B.get(), OUT.get(), numElements); + + try + { + ctx.synchronize(); + std::cerr << "Expected exception from out-of-range result but none was thrown!" << std::endl; + return EXIT_FAILURE; + } + catch (const std::domain_error& e) + { + std::cerr << "Caught expected error: " << e.what() << std::endl; + std::cerr << "Test PASSED\n"; + return 0; + } +} diff --git a/test/test_cuda_bounded_float_div.cu b/test/test_cuda_bounded_float_div.cu new file mode 100644 index 0000000..16a9ad2 --- /dev/null +++ b/test/test_cuda_bounded_float_div.cu @@ -0,0 +1,63 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" + +#include + +using test_type = boost::safe_numbers::bounded_float<-1000000.0f, 1000000.0f>; +using basis_type = test_type::basis_type; + +__global__ void cuda_test(const test_type *a, const test_type *b, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + out[i] = a[i] / b[i]; + } +} + +int main() +{ + int numElements = 2048; + + cuda_managed_ptr A(numElements); + cuda_managed_ptr B(numElements); + cuda_managed_ptr OUT(numElements); + + std::mt19937_64 rng{13}; + std::uniform_real_distribution da{-500000.0f, 500000.0f}; + std::uniform_real_distribution db{1.0f, 1000.0f}; + for (int i = 0; i < numElements; ++i) + { + A[i] = test_type{basis_type{da(rng)}}; + B[i] = test_type{basis_type{db(rng)}}; + } + + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + + boost::safe_numbers::device_error_context ctx; + cuda_test<<>>(A.get(), B.get(), OUT.get(), numElements); + ctx.synchronize(); + + for (int i = 0; i < numElements; ++i) + { + if (OUT[i] != A[i] / B[i]) + { + std::cerr << "Result verification failed at element " << i << "!" << std::endl; + return EXIT_FAILURE; + } + } + + std::cout << "Test PASSED\n"; + return 0; +} diff --git a/test/test_cuda_bounded_float_mul.cu b/test/test_cuda_bounded_float_mul.cu new file mode 100644 index 0000000..b054088 --- /dev/null +++ b/test/test_cuda_bounded_float_mul.cu @@ -0,0 +1,61 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" + +#include + +using test_type = boost::safe_numbers::bounded_float<0.0f, 1000000.0f>; +using basis_type = test_type::basis_type; + +__global__ void cuda_test(const test_type *a, const test_type *b, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + out[i] = a[i] * b[i]; + } +} + +int main() +{ + int numElements = 2048; + + cuda_managed_ptr A(numElements); + cuda_managed_ptr B(numElements); + cuda_managed_ptr OUT(numElements); + + std::mt19937_64 rng{7}; + std::uniform_real_distribution dist{0.0f, 900.0f}; + for (int i = 0; i < numElements; ++i) + { + A[i] = test_type{basis_type{dist(rng)}}; + B[i] = test_type{basis_type{dist(rng)}}; + } + + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + + boost::safe_numbers::device_error_context ctx; + cuda_test<<>>(A.get(), B.get(), OUT.get(), numElements); + ctx.synchronize(); + + for (int i = 0; i < numElements; ++i) + { + if (OUT[i] != A[i] * B[i]) + { + std::cerr << "Result verification failed at element " << i << "!" << std::endl; + return EXIT_FAILURE; + } + } + + std::cout << "Test PASSED\n"; + return 0; +} diff --git a/test/test_cuda_bounded_float_sub.cu b/test/test_cuda_bounded_float_sub.cu new file mode 100644 index 0000000..1a5fba1 --- /dev/null +++ b/test/test_cuda_bounded_float_sub.cu @@ -0,0 +1,62 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" + +#include + +using test_type = boost::safe_numbers::bounded_float<-1000000.0f, 1000000.0f>; +using basis_type = test_type::basis_type; + +__global__ void cuda_test(const test_type *a, const test_type *b, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + out[i] = a[i] - b[i]; + } +} + +int main() +{ + int numElements = 2048; + + cuda_managed_ptr A(numElements); + cuda_managed_ptr B(numElements); + cuda_managed_ptr OUT(numElements); + + std::mt19937_64 rng{11}; + std::uniform_real_distribution dist{0.0f, 400000.0f}; + for (int i = 0; i < numElements; ++i) + { + A[i] = test_type{basis_type{dist(rng)}}; + B[i] = test_type{basis_type{dist(rng)}}; + } + + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + + boost::safe_numbers::device_error_context ctx; + cuda_test<<>>(A.get(), B.get(), OUT.get(), numElements); + ctx.synchronize(); + + for (int i = 0; i < numElements; ++i) + { + if (OUT[i] != A[i] - B[i]) + { + std::cerr << "Result verification failed at element " << i << "!" << std::endl; + return EXIT_FAILURE; + } + } + + std::cout << "Test PASSED\n"; + return 0; +} diff --git a/test/test_cuda_bounded_int_add.cu b/test/test_cuda_bounded_int_add.cu new file mode 100644 index 0000000..930a11b --- /dev/null +++ b/test/test_cuda_bounded_int_add.cu @@ -0,0 +1,61 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" + +#include + +using test_type = boost::safe_numbers::bounded_int<-1000000, 1000000>; + +__global__ void cuda_test(const test_type *a, const test_type *b, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + out[i] = a[i] + b[i]; + } +} + +int main() +{ + int numElements = 2048; + + cuda_managed_ptr A(numElements); + cuda_managed_ptr B(numElements); + cuda_managed_ptr OUT(numElements); + + std::mt19937_64 rng{42}; + std::uniform_int_distribution dist{-400000, 400000}; + for (int i = 0; i < numElements; ++i) + { + A[i] = test_type{dist(rng)}; + B[i] = test_type{dist(rng)}; + } + + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + + boost::safe_numbers::device_error_context ctx; + cuda_test<<>>(A.get(), B.get(), OUT.get(), numElements); + ctx.synchronize(); + + for (int i = 0; i < numElements; ++i) + { + if (OUT[i] != A[i] + B[i]) + { + std::cerr << "Result verification failed at element " << i << "!" << std::endl; + return EXIT_FAILURE; + } + } + + std::cout << "Test PASSED\n"; + return 0; +} diff --git a/test/test_cuda_bounded_int_add_error.cu b/test/test_cuda_bounded_int_add_error.cu new file mode 100644 index 0000000..db3a7a9 --- /dev/null +++ b/test/test_cuda_bounded_int_add_error.cu @@ -0,0 +1,59 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" + +#include + +using test_type = boost::safe_numbers::bounded_int<-1000, 1000>; + +__global__ void cuda_test(const test_type *a, const test_type *b, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + out[i] = a[i] + b[i]; + } +} + +int main() +{ + int numElements = 1024; + + cuda_managed_ptr A(numElements); + cuda_managed_ptr B(numElements); + cuda_managed_ptr OUT(numElements); + + // 600 + 600 = 1200 exceeds the [-1000, 1000] bound -> out of range + for (int i = 0; i < numElements; ++i) + { + A[i] = test_type{600}; + B[i] = test_type{600}; + } + + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + + boost::safe_numbers::device_error_context ctx; + cuda_test<<>>(A.get(), B.get(), OUT.get(), numElements); + + try + { + ctx.synchronize(); + std::cerr << "Expected exception from out-of-range result but none was thrown!" << std::endl; + return EXIT_FAILURE; + } + catch (const std::domain_error& e) + { + std::cerr << "Caught expected error: " << e.what() << std::endl; + std::cerr << "Test PASSED\n"; + return 0; + } +} diff --git a/test/test_cuda_bounded_int_div.cu b/test/test_cuda_bounded_int_div.cu new file mode 100644 index 0000000..a5cfa16 --- /dev/null +++ b/test/test_cuda_bounded_int_div.cu @@ -0,0 +1,62 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" + +#include + +using test_type = boost::safe_numbers::bounded_int<-1000000, 1000000>; + +__global__ void cuda_test(const test_type *a, const test_type *b, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + out[i] = a[i] / b[i]; + } +} + +int main() +{ + int numElements = 2048; + + cuda_managed_ptr A(numElements); + cuda_managed_ptr B(numElements); + cuda_managed_ptr OUT(numElements); + + std::mt19937_64 rng{13}; + std::uniform_int_distribution da{-1000000, 1000000}; + std::uniform_int_distribution db{1, 1000}; + for (int i = 0; i < numElements; ++i) + { + A[i] = test_type{da(rng)}; + B[i] = test_type{db(rng)}; + } + + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + + boost::safe_numbers::device_error_context ctx; + cuda_test<<>>(A.get(), B.get(), OUT.get(), numElements); + ctx.synchronize(); + + for (int i = 0; i < numElements; ++i) + { + if (OUT[i] != A[i] / B[i]) + { + std::cerr << "Result verification failed at element " << i << "!" << std::endl; + return EXIT_FAILURE; + } + } + + std::cout << "Test PASSED\n"; + return 0; +} diff --git a/test/test_cuda_bounded_int_mul.cu b/test/test_cuda_bounded_int_mul.cu new file mode 100644 index 0000000..19d058f --- /dev/null +++ b/test/test_cuda_bounded_int_mul.cu @@ -0,0 +1,61 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" + +#include + +using test_type = boost::safe_numbers::bounded_int<-1000000, 1000000>; + +__global__ void cuda_test(const test_type *a, const test_type *b, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + out[i] = a[i] * b[i]; + } +} + +int main() +{ + int numElements = 2048; + + cuda_managed_ptr A(numElements); + cuda_managed_ptr B(numElements); + cuda_managed_ptr OUT(numElements); + + std::mt19937_64 rng{7}; + std::uniform_int_distribution dist{-900, 900}; + for (int i = 0; i < numElements; ++i) + { + A[i] = test_type{dist(rng)}; + B[i] = test_type{dist(rng)}; + } + + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + + boost::safe_numbers::device_error_context ctx; + cuda_test<<>>(A.get(), B.get(), OUT.get(), numElements); + ctx.synchronize(); + + for (int i = 0; i < numElements; ++i) + { + if (OUT[i] != A[i] * B[i]) + { + std::cerr << "Result verification failed at element " << i << "!" << std::endl; + return EXIT_FAILURE; + } + } + + std::cout << "Test PASSED\n"; + return 0; +} diff --git a/test/test_cuda_bounded_int_sub.cu b/test/test_cuda_bounded_int_sub.cu new file mode 100644 index 0000000..926389c --- /dev/null +++ b/test/test_cuda_bounded_int_sub.cu @@ -0,0 +1,61 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" + +#include + +using test_type = boost::safe_numbers::bounded_int<-1000000, 1000000>; + +__global__ void cuda_test(const test_type *a, const test_type *b, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + out[i] = a[i] - b[i]; + } +} + +int main() +{ + int numElements = 2048; + + cuda_managed_ptr A(numElements); + cuda_managed_ptr B(numElements); + cuda_managed_ptr OUT(numElements); + + std::mt19937_64 rng{11}; + std::uniform_int_distribution dist{-500000, 500000}; + for (int i = 0; i < numElements; ++i) + { + A[i] = test_type{dist(rng)}; + B[i] = test_type{dist(rng)}; + } + + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + + boost::safe_numbers::device_error_context ctx; + cuda_test<<>>(A.get(), B.get(), OUT.get(), numElements); + ctx.synchronize(); + + for (int i = 0; i < numElements; ++i) + { + if (OUT[i] != A[i] - B[i]) + { + std::cerr << "Result verification failed at element " << i << "!" << std::endl; + return EXIT_FAILURE; + } + } + + std::cout << "Test PASSED\n"; + return 0; +} diff --git a/test/test_cuda_bounded_ops.cu b/test/test_cuda_bounded_ops.cu new file mode 100644 index 0000000..ed5261e --- /dev/null +++ b/test/test_cuda_bounded_ops.cu @@ -0,0 +1,83 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://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. + +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" + +#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>; + +__global__ void ops_kernel(int* ok) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i == 0) + { + // bounded_uint: modulo, increment, decrement, compound assignment + const bu a {100}; + const bu b {30}; + bu inc {a}; + ++inc; + inc--; + bu comp {a}; + comp += b; + comp -= b; + comp *= bu{2}; + comp /= bu{2}; + const bool r1 {(a % b == bu{10}) && (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 q {bf::basis_type{3.0f}}; + bf pc {p}; + pc *= q; + pc /= q; + const bool r3 {(pc == p) && (p < q)}; + + *ok = (r1 && r2 && r3) ? 1 : 0; + } +} + +int main() +{ + int* ok {nullptr}; + cudaMallocManaged(&ok, sizeof(int)); + cudaDeviceSynchronize(); + *ok = 0; + + boost::safe_numbers::device_error_context ctx; + ops_kernel<<<1, 1>>>(ok); + ctx.synchronize(); + + if (*ok == 1) + { + std::cout << "Test PASSED\n"; + cudaFree(ok); + return 0; + } + + std::cerr << "Result verification failed!" << std::endl; + cudaFree(ok); + return EXIT_FAILURE; +} diff --git a/test/test_cuda_bounded_uint_add.cu b/test/test_cuda_bounded_uint_add.cu new file mode 100644 index 0000000..0c32f0c --- /dev/null +++ b/test/test_cuda_bounded_uint_add.cu @@ -0,0 +1,61 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" + +#include + +using test_type = boost::safe_numbers::bounded_uint<0u, 1000000u>; + +__global__ void cuda_test(const test_type *a, const test_type *b, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + out[i] = a[i] + b[i]; + } +} + +int main() +{ + int numElements = 2048; + + cuda_managed_ptr A(numElements); + cuda_managed_ptr B(numElements); + cuda_managed_ptr OUT(numElements); + + std::mt19937_64 rng{42}; + std::uniform_int_distribution dist{0, 400000}; + for (int i = 0; i < numElements; ++i) + { + A[i] = test_type{dist(rng)}; + B[i] = test_type{dist(rng)}; + } + + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + + boost::safe_numbers::device_error_context ctx; + cuda_test<<>>(A.get(), B.get(), OUT.get(), numElements); + ctx.synchronize(); + + for (int i = 0; i < numElements; ++i) + { + if (OUT[i] != A[i] + B[i]) + { + std::cerr << "Result verification failed at element " << i << "!" << std::endl; + return EXIT_FAILURE; + } + } + + std::cout << "Test PASSED\n"; + return 0; +} diff --git a/test/test_cuda_bounded_uint_add_error.cu b/test/test_cuda_bounded_uint_add_error.cu new file mode 100644 index 0000000..90108c6 --- /dev/null +++ b/test/test_cuda_bounded_uint_add_error.cu @@ -0,0 +1,59 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" + +#include + +using test_type = boost::safe_numbers::bounded_uint<0u, 1000u>; + +__global__ void cuda_test(const test_type *a, const test_type *b, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + out[i] = a[i] + b[i]; + } +} + +int main() +{ + int numElements = 1024; + + cuda_managed_ptr A(numElements); + cuda_managed_ptr B(numElements); + cuda_managed_ptr OUT(numElements); + + // 600 + 600 = 1200 exceeds the [0, 1000] bound -> out of range + for (int i = 0; i < numElements; ++i) + { + A[i] = test_type{600}; + B[i] = test_type{600}; + } + + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + + boost::safe_numbers::device_error_context ctx; + cuda_test<<>>(A.get(), B.get(), OUT.get(), numElements); + + try + { + ctx.synchronize(); + std::cerr << "Expected exception from out-of-range result but none was thrown!" << std::endl; + return EXIT_FAILURE; + } + catch (const std::domain_error& e) + { + std::cerr << "Caught expected error: " << e.what() << std::endl; + std::cerr << "Test PASSED\n"; + return 0; + } +} diff --git a/test/test_cuda_bounded_uint_div.cu b/test/test_cuda_bounded_uint_div.cu new file mode 100644 index 0000000..0f9fda2 --- /dev/null +++ b/test/test_cuda_bounded_uint_div.cu @@ -0,0 +1,62 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" + +#include + +using test_type = boost::safe_numbers::bounded_uint<0u, 1000000u>; + +__global__ void cuda_test(const test_type *a, const test_type *b, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + out[i] = a[i] / b[i]; + } +} + +int main() +{ + int numElements = 2048; + + cuda_managed_ptr A(numElements); + cuda_managed_ptr B(numElements); + cuda_managed_ptr OUT(numElements); + + std::mt19937_64 rng{13}; + std::uniform_int_distribution da{0, 1000000}; + std::uniform_int_distribution db{1, 1000}; + for (int i = 0; i < numElements; ++i) + { + A[i] = test_type{da(rng)}; + B[i] = test_type{db(rng)}; + } + + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + + boost::safe_numbers::device_error_context ctx; + cuda_test<<>>(A.get(), B.get(), OUT.get(), numElements); + ctx.synchronize(); + + for (int i = 0; i < numElements; ++i) + { + if (OUT[i] != A[i] / B[i]) + { + std::cerr << "Result verification failed at element " << i << "!" << std::endl; + return EXIT_FAILURE; + } + } + + std::cout << "Test PASSED\n"; + return 0; +} diff --git a/test/test_cuda_bounded_uint_mul.cu b/test/test_cuda_bounded_uint_mul.cu new file mode 100644 index 0000000..1d3ae22 --- /dev/null +++ b/test/test_cuda_bounded_uint_mul.cu @@ -0,0 +1,61 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" + +#include + +using test_type = boost::safe_numbers::bounded_uint<0u, 1000000u>; + +__global__ void cuda_test(const test_type *a, const test_type *b, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + out[i] = a[i] * b[i]; + } +} + +int main() +{ + int numElements = 2048; + + cuda_managed_ptr A(numElements); + cuda_managed_ptr B(numElements); + cuda_managed_ptr OUT(numElements); + + std::mt19937_64 rng{7}; + std::uniform_int_distribution dist{0, 900}; + for (int i = 0; i < numElements; ++i) + { + A[i] = test_type{dist(rng)}; + B[i] = test_type{dist(rng)}; + } + + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + + boost::safe_numbers::device_error_context ctx; + cuda_test<<>>(A.get(), B.get(), OUT.get(), numElements); + ctx.synchronize(); + + for (int i = 0; i < numElements; ++i) + { + if (OUT[i] != A[i] * B[i]) + { + std::cerr << "Result verification failed at element " << i << "!" << std::endl; + return EXIT_FAILURE; + } + } + + std::cout << "Test PASSED\n"; + return 0; +} diff --git a/test/test_cuda_bounded_uint_sub.cu b/test/test_cuda_bounded_uint_sub.cu new file mode 100644 index 0000000..61968fb --- /dev/null +++ b/test/test_cuda_bounded_uint_sub.cu @@ -0,0 +1,64 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +#include +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" + +#include + +using test_type = boost::safe_numbers::bounded_uint<0u, 1000000u>; + +__global__ void cuda_test(const test_type *a, const test_type *b, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + out[i] = a[i] - b[i]; + } +} + +int main() +{ + int numElements = 2048; + + cuda_managed_ptr A(numElements); + cuda_managed_ptr B(numElements); + cuda_managed_ptr OUT(numElements); + + std::mt19937_64 rng{11}; + std::uniform_int_distribution db{0, 500000}; + std::uniform_int_distribution dd{0, 500000}; + for (int i = 0; i < numElements; ++i) + { + const std::uint32_t b {db(rng)}; + const std::uint32_t a {static_cast(b + dd(rng))}; + A[i] = test_type{a}; + B[i] = test_type{b}; + } + + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + + boost::safe_numbers::device_error_context ctx; + cuda_test<<>>(A.get(), B.get(), OUT.get(), numElements); + ctx.synchronize(); + + for (int i = 0; i < numElements; ++i) + { + if (OUT[i] != A[i] - B[i]) + { + std::cerr << "Result verification failed at element " << i << "!" << std::endl; + return EXIT_FAILURE; + } + } + + std::cout << "Test PASSED\n"; + return 0; +} diff --git a/test/test_cuda_invalid_argument.cu b/test/test_cuda_invalid_argument.cu new file mode 100644 index 0000000..504937d --- /dev/null +++ b/test/test_cuda_invalid_argument.cu @@ -0,0 +1,72 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +// Verifies that a std::invalid_argument raised through the device error +// reporter is rethrown on the host as std::invalid_argument (and not folded +// into std::domain_error). This exercises both host-side reporting inside a +// CUDA translation unit and the device-kernel -> synchronize() round-trip. + +#include +#include +#include +#include +#include + +#include + +using namespace boost::safe_numbers::literals; + +// Records an invalid_argument on the device, exactly as +// BOOST_SAFE_NUMBERS_THROW_EXCEPTION(std::invalid_argument, ...) does. +__global__ void invalid_argument_kernel() +{ + if (blockDim.x * blockIdx.x + threadIdx.x == 0) + { + boost::safe_numbers::detail::report_device_error( + boost::safe_numbers::detail::to_exception_enum(), + __FILE__, __LINE__, "invalid_argument on device"); + } +} + +int main() +{ + int failures {0}; + + // Host reporting inside a CUDA TU: a bad string literal reports through the + // host branch of report_device_error and must surface as invalid_argument. + try + { + static_cast(operator""_u128("-1")); + std::cerr << "host: expected std::invalid_argument, none thrown" << std::endl; + ++failures; + } + catch (const std::invalid_argument& e) + { + std::cerr << "host literal: caught std::invalid_argument: " << e.what() << std::endl; + } + + // Device round-trip: the kernel records invalid_argument and synchronize() + // must rethrow it as std::invalid_argument. + boost::safe_numbers::device_error_context ctx; + invalid_argument_kernel<<<1, 1>>>(); + try + { + ctx.synchronize(); + std::cerr << "device: expected std::invalid_argument, none thrown" << std::endl; + ++failures; + } + catch (const std::invalid_argument& e) + { + std::cerr << "device kernel: caught std::invalid_argument: " << e.what() << std::endl; + } + + if (failures == 0) + { + std::cerr << "Test PASSED\n"; + return 0; + } + + return EXIT_FAILURE; +} diff --git a/test/test_device_error_mapping.cpp b/test/test_device_error_mapping.cpp new file mode 100644 index 0000000..2b52044 --- /dev/null +++ b/test/test_device_error_mapping.cpp @@ -0,0 +1,26 @@ +// Copyright Matt Borland 2026. +// Use, modification and distribution are subject to the +// Boost Software License, Version 1.0. (See accompanying file +// LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt) + +// Guards the std-exception-type to error_category mapping used by the CUDA +// device error reporter. In particular std::invalid_argument must map to its +// 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 + +using boost::safe_numbers::detail::exception_type; +using boost::safe_numbers::detail::to_exception_enum; + +static_assert(to_exception_enum() == exception_type::domain_error); +static_assert(to_exception_enum() == exception_type::overflow); +static_assert(to_exception_enum() == exception_type::underflow); +static_assert(to_exception_enum() == exception_type::invalid_argument); +static_assert(to_exception_enum() == exception_type::unknown); + +int main() +{ + return 0; +}