Implement and test bounded types on CUDA

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]

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 <<cuda_floats>>.
 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`
 |===

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
 

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 <iostream>
+#include <stdexcept>
+#include <boost/safe_numbers/bounded_integers.hpp>
+#include <boost/safe_numbers/bounded_floats.hpp>
+#include <boost/safe_numbers/cuda_error_reporting.hpp>
+
+#include <cuda_runtime.h>
+
+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 <typename T>
+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>()};
+    bu* ub {managed<bu>()};
+    bu* uo {managed<bu>()};
+    *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>()};
+    bi* ib {managed<bi>()};
+    bi* io {managed<bi>()};
+    *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>()};
+    bf* fb {managed<bf>()};
+    bf* fo {managed<bf>()};
+    *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;
+}

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<underlying_type>(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 <typename OtherBasis>
         requires (detail::is_compatible_float_type<OtherBasis>)
-    [[nodiscard]] explicit constexpr operator OtherBasis() const
+    BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] explicit constexpr operator OtherBasis() const
     {
         const auto raw {static_cast<underlying_type>(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,28 +128,28 @@ class bounded_float
     }
 
     template <auto Min2, auto Max2>
-    [[nodiscard]] explicit constexpr operator bounded_float<Min2, Max2>() const
+    BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] explicit constexpr operator bounded_float<Min2, Max2>() const
     {
         using target_basis = typename bounded_float<Min2, Max2>::basis_type;
         using target_underlying = detail::underlying_type_t<target_basis>;
         const auto raw {static_cast<underlying_type>(basis_)};
         return bounded_float<Min2, Max2>{target_basis{static_cast<target_underlying>(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<Min, Max> rhs) -> bounded_float&;
+    BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator+=(bounded_float<Min, Max> rhs) -> bounded_float&;
 
-    constexpr auto operator-=(bounded_float<Min, Max> rhs) -> bounded_float&;
+    BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator-=(bounded_float<Min, Max> rhs) -> bounded_float&;
 
-    constexpr auto operator*=(bounded_float<Min, Max> rhs) -> bounded_float&;
+    BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator*=(bounded_float<Min, Max> rhs) -> bounded_float&;
 
-    constexpr auto operator/=(bounded_float<Min, Max> rhs) -> bounded_float&;
+    BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto operator/=(bounded_float<Min, Max> rhs) -> bounded_float&;
 };
 
 // ------------------------------
@@ -155,7 +161,7 @@ class bounded_float
 // The bounded_float constructor then re-validates the result against [Min, Max].
 
 template <auto Min, auto Max>
-[[nodiscard]] constexpr auto operator+(const bounded_float<Min, Max> lhs,
+BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator+(const bounded_float<Min, Max> lhs,
                                        const bounded_float<Min, Max> rhs) -> bounded_float<Min, Max>
 {
     using basis = typename bounded_float<Min, Max>::basis_type;
@@ -166,7 +172,7 @@ template <auto Min, auto Max>
 }
 
 template <auto Min, auto Max>
-[[nodiscard]] constexpr auto operator-(const bounded_float<Min, Max> lhs,
+BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator-(const bounded_float<Min, Max> lhs,
                                        const bounded_float<Min, Max> rhs) -> bounded_float<Min, Max>
 {
     using basis = typename bounded_float<Min, Max>::basis_type;
@@ -177,7 +183,7 @@ template <auto Min, auto Max>
 }
 
 template <auto Min, auto Max>
-[[nodiscard]] constexpr auto operator*(const bounded_float<Min, Max> lhs,
+BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator*(const bounded_float<Min, Max> lhs,
                                        const bounded_float<Min, Max> rhs) -> bounded_float<Min, Max>
 {
     using basis = typename bounded_float<Min, Max>::basis_type;
@@ -188,7 +194,7 @@ template <auto Min, auto Max>
 }
 
 template <auto Min, auto Max>
-[[nodiscard]] constexpr auto operator/(const bounded_float<Min, Max> lhs,
+BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto operator/(const bounded_float<Min, Max> lhs,
                                        const bounded_float<Min, Max> rhs) -> bounded_float<Min, Max>
 {
     using basis = typename bounded_float<Min, Max>::basis_type;
@@ -209,7 +215,7 @@ template <auto Min, auto Max>
               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<Min, Max>::operator+=(bounded_float<Min, Max> rhs) -> bounded_float&
+BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_float<Min, Max>::operator+=(bounded_float<Min, Max> rhs) -> bounded_float&
 {
     *this = *this + rhs;
     return *this;
@@ -222,7 +228,7 @@ template <auto Min, auto Max>
               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<Min, Max>::operator-=(bounded_float<Min, Max> rhs) -> bounded_float&
+BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_float<Min, Max>::operator-=(bounded_float<Min, Max> rhs) -> bounded_float&
 {
     *this = *this - rhs;
     return *this;
@@ -235,7 +241,7 @@ template <auto Min, auto Max>
               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<Min, Max>::operator*=(bounded_float<Min, Max> rhs) -> bounded_float&
+BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_float<Min, Max>::operator*=(bounded_float<Min, Max> rhs) -> bounded_float&
 {
     *this = *this * rhs;
     return *this;
@@ -248,7 +254,7 @@ template <auto Min, auto Max>
               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<Min, Max>::operator/=(bounded_float<Min, Max> rhs) -> bounded_float&
+BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto bounded_float<Min, Max>::operator/=(bounded_float<Min, Max> rhs) -> bounded_float&
 {
     *this = *this / rhs;
     return *this;