diff --git a/doc/modules/ROOT/pages/benchmarks.adoc b/doc/modules/ROOT/pages/benchmarks.adoc index 61452ab..0200a9a 100644 --- a/doc/modules/ROOT/pages/benchmarks.adoc +++ b/doc/modules/ROOT/pages/benchmarks.adoc @@ -86,6 +86,9 @@ That is deliberate: wrapping integer arithmetic is associative, so the optimizer Division and modulo cannot be reassociated, and interleaving them forces every step to execute for both contenders. The floating point chains are per operation (eight operations plus one restore step of the inverse kind), because IEEE arithmetic is not associative and cannot be folded. +Every family also measures the non-throwing forms against the same raw baseline: a `sat_add` streaming row (xref:unsigned_integers.adoc[`saturating_add`]), `mixed (sat)` and `mixed (ovf)` integer chains built from the `saturating_*` and `overflowing_*` functions, and `(ovf)` float chains built from the xref:floats.adoc#floats_overflowing_arithmetic[deferred float family] with the flags accumulated and folded into the result. +These rows show what checked arithmetic costs when the error report is deferred to a boundary instead of raised per operation. + Every timed region is verified afterwards: the safe results must match the builtin results bit for bit across all elements, a slice is recomputed on the host, and the process exit code reflects any mismatch. Both benchmarks accept optional positional arguments `[stream_elements] [chain_cycles] [launches] [chain_threads]` for scaling the run to slower or faster devices. @@ -125,59 +128,60 @@ NOTE: `icpx` compiles with `-ffp-model=fast` by default, which grants the optimi Run on an NVIDIA GeForce RTX 3060 (sm_86, 28 SMs) with CUDA 13.1 on Ubuntu 24.04, compiled at `-O2` with `--expt-relaxed-constexpr`, using the default sizes. .Streaming, ratio of safe time to basis type time (RTX 3060) -[cols="1,>1,>1,>1,>1,>1",options="header"] -|=== -| Type | Addition | Subtraction | Multiplication | Division | Modulo -| `u8` | 1.05 | 1.04 | 1.05 | 1.03 | 1.02 -| `u16` | 1.01 | 1.00 | 1.00 | 1.00 | 1.00 -| `u32` | 1.00 | 1.00 | 1.00 | 1.00 | 1.00 -| `u64` | 1.00 | 1.00 | 1.00 | 1.00 | 1.00 -| `u128` | 1.00 | 1.00 | 6.78 | 1.00 | 1.00 -| `i8` | 1.23 | 1.23 | 1.22 | 1.16 | 1.16 -| `i16` | 1.05 | 1.05 | 1.05 | 1.04 | 1.04 -| `i32` | 1.00 | 1.00 | 1.00 | 1.00 | 1.00 -| `i64` | 1.00 | 1.00 | 1.00 | 1.00 | 1.00 -| `i128` | 1.00 | 1.00 | 6.78 | 1.00 | 1.00 -| `f32` | 1.00 | 1.00 | 1.00 | 1.00 | n/a -| `f64` | 1.00 | 1.00 | 1.00 | 1.44 | n/a +[cols="1,>1,>1,>1,>1,>1,>1",options="header"] +|=== +| Type | Addition | Subtraction | Multiplication | Division | Modulo | Saturating addition +| `u8` | 1.05 | 1.04 | 1.05 | 1.03 | 1.02 | 1.02 +| `u16` | 1.00 | 1.00 | 1.00 | 1.00 | 1.00 | 0.99 +| `u32` | 1.00 | 1.00 | 1.00 | 1.00 | 1.00 | 1.00 +| `u64` | 1.00 | 1.00 | 1.00 | 1.00 | 1.00 | 1.00 +| `u128` | 1.00 | 1.00 | 6.77 | 1.00 | 1.00 | 1.00 +| `i8` | 1.23 | 1.22 | 1.22 | 1.16 | 1.15 | 1.04 +| `i16` | 1.05 | 1.05 | 1.04 | 1.04 | 1.03 | 1.00 +| `i32` | 1.00 | 1.00 | 1.00 | 1.00 | 1.00 | 1.00 +| `i64` | 1.00 | 1.00 | 1.00 | 1.00 | 1.00 | 1.00 +| `i128` | 1.00 | 1.00 | 6.79 | 1.00 | 1.00 | 1.00 +| `f32` | 1.00 | 1.00 | 1.00 | 1.00 | n/a | n/a +| `f64` | 1.00 | 1.00 | 1.00 | 1.44 | n/a | n/a |=== In the streaming regime the checks are essentially free: the kernels are limited by memory bandwidth and the check instructions execute in the shadow of the loads. -The only rows above parity are the 8-bit types (up to 1.23x, since a byte of traffic buys less time to hide the check), `f64` division (1.44x), and 128-bit multiplication (6.78x). +The only rows above parity are the 8-bit types (up to 1.23x, since a byte of traffic buys less time to hide the check), `f64` division (1.44x), and 128-bit multiplication (about 6.8x). The multiplication overflow check on the 128-bit types is the one genuinely expensive check in the library, since it must be computed from 64-bit limbs. -.Compute-bound integer chain (RTX 3060) -[cols="1,>1,>1,>1",options="header"] +.Compute-bound integer chain (RTX 3060): throwing operators and non-throwing forms against the raw chain +[cols="1,>1,>1,>1,>1,>1,>1,>1",options="header"] +|=== +| Type | Basis type (ns/op) | Throwing (ns/op) | Ratio | Saturating (ns/op) | Ratio | Overflowing (ns/op) | Ratio +| `u8` | 0.0030 | 0.0039 | 1.30 | 0.0037 | 1.24 | 0.0036 | 1.21 +| `u16` | 0.0030 | 0.0039 | 1.30 | 0.0037 | 1.23 | 0.0037 | 1.23 +| `u32` | 0.0012 | 0.0034 | 2.82 | 0.0030 | 2.55 | 0.0030 | 2.50 +| `u64` | 0.0031 | 0.0045 | 1.46 | 0.0044 | 1.44 | 0.0043 | 1.38 +| `u128` | 0.0094 | 0.0964 | 10.30 | 0.0884 | 9.43 | 0.0925 | 10.28 +| `i8` | 0.0020 | 0.0054 | 2.68 | 0.0056 | 2.78 | 0.0052 | 2.61 +| `i16` | 0.0020 | 0.0058 | 2.93 | 0.0058 | 2.96 | 0.0055 | 2.80 +| `i32` | 0.0018 | 0.0051 | 2.85 | 0.0052 | 2.92 | 0.0051 | 2.83 +| `i64` | 0.0031 | 0.0068 | 2.22 | 0.0071 | 2.31 | 0.0067 | 2.19 +| `i128` | 0.0121 | 0.0940 | 7.78 | 0.0921 | 7.62 | 0.0919 | 7.92 +|=== + +.Compute-bound floating point chains (RTX 3060): throwing operators and the deferred (overflowing) family +[cols="1,1,>1,>1,>1,>1,>1",options="header"] |=== -| Type | Basis type (ns/op) | Boost.SafeNumbers (ns/op) | Ratio -| `u8` | 0.0030 | 0.0039 | 1.30 -| `u16` | 0.0030 | 0.0039 | 1.30 -| `u32` | 0.0012 | 0.0034 | 2.82 -| `u64` | 0.0031 | 0.0045 | 1.46 -| `u128` | 0.0094 | 0.0963 | 10.30 -| `i8` | 0.0021 | 0.0056 | 2.68 -| `i16` | 0.0020 | 0.0058 | 2.92 -| `i32` | 0.0018 | 0.0051 | 2.85 -| `i64` | 0.0031 | 0.0068 | 2.22 -| `i128` | 0.0121 | 0.0944 | 7.81 -|=== - -.Compute-bound floating point chains (RTX 3060) -[cols="1,1,>1,>1,>1",options="header"] -|=== -| Type | Operation | Basis type (ns/op) | Boost.SafeNumbers (ns/op) | Ratio -| `f32` | Addition | 0.0002 | 0.0037 | 17.44 -| `f32` | Subtraction | 0.0002 | 0.0037 | 17.74 -| `f32` | Multiplication | 0.0004 | 0.0038 | 8.73 -| `f32` | Division | 0.0026 | 0.0057 | 2.22 -| `f64` | Addition | 0.0124 | 0.0494 | 4.00 -| `f64` | Subtraction | 0.0133 | 0.0489 | 3.68 -| `f64` | Multiplication | 0.0200 | 0.0561 | 2.80 -| `f64` | Division | 0.0770 | 0.1151 | 1.50 +| Type | Operation | Basis type (ns/op) | Throwing (ns/op) | Ratio | Overflowing (ns/op) | Ratio +| `f32` | Addition | 0.0002 | 0.0037 | 18.31 | 0.0011 | 5.39 +| `f32` | Subtraction | 0.0002 | 0.0037 | 17.66 | 0.0011 | 5.37 +| `f32` | Multiplication | 0.0004 | 0.0039 | 8.71 | 0.0012 | 2.69 +| `f32` | Division | 0.0026 | 0.0057 | 2.22 | 0.0027 | 1.05 +| `f64` | Addition | 0.0124 | 0.0495 | 4.00 | 0.0501 | 4.21 +| `f64` | Subtraction | 0.0124 | 0.0495 | 4.01 | 0.0501 | 4.21 +| `f64` | Multiplication | 0.0209 | 0.0579 | 2.77 | 0.0593 | 2.94 +| `f64` | Division | 0.0779 | 0.1157 | 1.49 | 0.1196 | 1.55 |=== With the arithmetic pinned in registers the checked integer types cost 1.3x to 2.9x through 64 bits, and the 128-bit chains are dominated by the multiplication check (7.8x to 10.3x). -For the floating point chains the baseline operation is a single hardware instruction, so the fixed cost of the checks multiplies it: up to 17.7x for `f32`, and 1.5x to 4.0x for `f64`, whose baseline instruction is already slow on a consumer GPU. +For the floating point chains the baseline operation is a single hardware instruction, so the fixed cost of the checks multiplies it: up to 18.3x for `f32`, and 1.5x to 4.0x for `f64`, whose baseline instruction is already slow on a consumer GPU. +On a GPU the non-throwing integer forms buy little, because the throwing operators already compile to a well predicted branch; the deferred float family is the exception, cutting the `f32` chains from 18.3x to 5.4x and `f32` division to parity by dropping the per-operation classification. Real kernels sit between the two regimes, and usually much closer to the streaming one. [#sycl_cpu_benchmarks] @@ -187,59 +191,63 @@ Run on the OpenCL CPU device of an Intel Core i9-11900K (8 cores) with oneAPI ic A CPU device executes a kernel by vectorizing it across work items, and that shapes these results: a branch-free kernel body runs 8 or 16 lanes wide, and the branches of the safety checks inhibit exactly that transformation. .Streaming, ratio of safe time to basis type time (SYCL CPU device) -[cols="1,>1,>1,>1,>1,>1",options="header"] -|=== -| Type | Addition | Subtraction | Multiplication | Division | Modulo -| `u8` | 5.86 | 6.59 | 14.80 | 1.08 | 1.01 -| `u16` | 1.05 | 1.07 | 1.78 | 1.27 | 1.28 -| `u32` | 1.04 | 1.05 | 1.04 | 1.02 | 1.00 -| `u64` | 1.03 | 1.03 | 1.03 | 1.00 | 1.00 -| `u128` | 1.00 | 1.01 | 1.05 | 1.01 | 0.99 -| `i8` | 9.48 | 9.61 | 8.75 | 1.15 | 1.16 -| `i16` | 1.58 | 1.46 | 1.60 | 1.75 | 1.64 -| `i32` | 1.05 | 1.05 | 1.05 | 1.03 | 1.07 -| `i64` | 1.04 | 1.03 | 1.03 | 1.01 | 1.01 -| `i128` | 1.00 | 1.01 | 1.11 | 1.00 | 1.00 -| `f32` | 1.37 | 1.36 | 1.38 | 1.72 | n/a -| `f64` | 1.06 | 1.05 | 1.05 | 1.14 | n/a +[cols="1,>1,>1,>1,>1,>1,>1",options="header"] +|=== +| Type | Addition | Subtraction | Multiplication | Division | Modulo | Saturating addition +| `u8` | 5.47 | 6.02 | 11.60 | 1.01 | 1.01 | 0.97 +| `u16` | 1.03 | 1.06 | 1.87 | 1.15 | 1.25 | 1.02 +| `u32` | 1.03 | 1.04 | 1.04 | 1.02 | 1.00 | 1.00 +| `u64` | 1.03 | 1.03 | 1.04 | 1.00 | 1.00 | 1.00 +| `u128` | 1.00 | 1.01 | 1.05 | 1.01 | 0.99 | 0.99 +| `i8` | 9.69 | 10.84 | 11.32 | 1.14 | 1.14 | 1.52 +| `i16` | 1.50 | 1.50 | 1.44 | 1.59 | 1.59 | 1.07 +| `i32` | 1.00 | 1.00 | 1.00 | 0.98 | 0.99 | 0.96 +| `i64` | 1.04 | 1.04 | 1.03 | 1.01 | 1.01 | 1.00 +| `i128` | 1.01 | 1.01 | 1.11 | 1.01 | 1.00 | 0.99 +| `f32` | 1.32 | 1.35 | 1.35 | 1.74 | n/a | n/a +| `f64` | 1.05 | 1.05 | 1.05 | 1.12 | n/a | n/a |=== For 32 bits and up the streaming results match the CUDA picture: the arrays no longer fit in cache, the kernels are bound by memory, and the checks cost 0 to 7 percent. -The 8-bit rows stand out (up to 14.8x) because four million 1-byte elements fit inside the processor's cache: those kernels are compute bound, so they measure the vectorization gap rather than memory bandwidth. - -.Compute-bound integer chain (SYCL CPU device) -[cols="1,>1,>1,>1",options="header"] -|=== -| Type | Basis type (ns/op) | Boost.SafeNumbers (ns/op) | Ratio -| `u8` | 0.0810 | 0.1229 | 1.52 -| `u16` | 0.0845 | 0.1486 | 1.76 -| `u32` | 0.0220 | 0.0833 | 3.79 -| `u64` | 0.1263 | 0.1431 | 1.13 -| `u128` | 0.3279 | 0.5028 | 1.53 -| `i8` | 0.1300 | 0.2062 | 1.59 -| `i16` | 0.0243 | 0.2057 | 8.48 -| `i32` | 0.0215 | 0.2043 | 9.48 -| `i64` | 0.1266 | 0.1838 | 1.45 -| `i128` | 0.4155 | 0.6179 | 1.49 -|=== - -.Compute-bound floating point chains (SYCL CPU device) -[cols="1,1,>1,>1,>1",options="header"] -|=== -| Type | Operation | Basis type (ns/op) | Boost.SafeNumbers (ns/op) | Ratio -| `f32` | Addition | 0.0035 | 0.6207 | 175.43 -| `f32` | Subtraction | 0.0035 | 0.6072 | 171.47 -| `f32` | Multiplication | 0.0050 | 0.6302 | 126.55 -| `f32` | Division | 0.0149 | 0.7543 | 50.64 -| `f64` | Addition | 0.0036 | 0.6496 | 181.82 -| `f64` | Subtraction | 0.0036 | 0.6516 | 182.26 -| `f64` | Multiplication | 0.0063 | 0.6594 | 104.15 -| `f64` | Division | 0.0451 | 0.7733 | 17.16 +The 8-bit rows stand out (up to 11.6x) because four million 1-byte elements fit inside the processor's cache: those kernels are compute bound, so they measure the vectorization gap rather than memory bandwidth. +The saturating addition column shows the escape hatch: `saturating_add` compiles to the native saturating vector instructions, so even the 8-bit rows return to parity. + +.Compute-bound integer chain (SYCL CPU device): throwing operators and non-throwing forms against the raw chain +[cols="1,>1,>1,>1,>1,>1,>1,>1",options="header"] +|=== +| Type | Basis type (ns/op) | Throwing (ns/op) | Ratio | Saturating (ns/op) | Ratio | Overflowing (ns/op) | Ratio +| `u8` | 0.0803 | 0.1198 | 1.49 | 0.1185 | 1.51 | 0.1149 | 1.43 +| `u16` | 0.0833 | 0.1567 | 1.88 | 0.1224 | 1.46 | 0.1434 | 1.72 +| `u32` | 0.0220 | 0.0941 | 4.27 | 0.0742 | 3.34 | 0.0697 | 3.16 +| `u64` | 0.1259 | 0.1317 | 1.05 | 0.1248 | 0.99 | 0.1271 | 1.00 +| `u128` | 0.3259 | 0.5082 | 1.56 | 0.4989 | 1.55 | 0.5189 | 1.59 +| `i8` | 0.1307 | 0.2153 | 1.65 | 0.1344 | 1.06 | 0.1459 | 1.11 +| `i16` | 0.0233 | 0.2057 | 8.82 | 0.1346 | 5.76 | 0.1406 | 5.90 +| `i32` | 0.0215 | 0.2025 | 9.42 | 0.1332 | 6.05 | 0.1415 | 6.46 +| `i64` | 0.1270 | 0.1858 | 1.46 | 0.1473 | 1.16 | 0.1485 | 1.17 +| `i128` | 0.4167 | 0.6075 | 1.46 | 0.8602 | 2.06 | 0.9172 | 2.22 +|=== + +.Compute-bound floating point chains (SYCL CPU device): throwing operators and the deferred (overflowing) family +[cols="1,1,>1,>1,>1,>1,>1",options="header"] +|=== +| Type | Operation | Basis type (ns/op) | Throwing (ns/op) | Ratio | Overflowing (ns/op) | Ratio +| `f32` | Addition | 0.0035 | 0.6126 | 173.03 | 0.0060 | 1.68 +| `f32` | Subtraction | 0.0035 | 0.6066 | 171.30 | 0.0072 | 2.11 +| `f32` | Multiplication | 0.0049 | 0.6275 | 127.59 | 0.0071 | 1.43 +| `f32` | Division | 0.0150 | 0.7521 | 50.01 | 0.0163 | 1.13 +| `f64` | Addition | 0.0036 | 0.6457 | 179.70 | 0.0098 | 2.73 +| `f64` | Subtraction | 0.0036 | 0.8446 | 234.65 | 0.0108 | 2.85 +| `f64` | Multiplication | 0.0064 | 0.6579 | 102.24 | 0.0135 | 2.05 +| `f64` | Division | 0.0450 | 0.7644 | 17.00 | 0.0475 | 1.05 |=== In the chains the dominant cost on a CPU device is lost vectorization rather than the check instructions themselves. -Where the baseline cannot vectorize either (the 64-bit and 128-bit chains, and the 8-bit chains whose narrow division defeats the vectorizer), the checked types cost 1.1x to 1.8x, in line with the host benchmarks. -Where the baseline vectorizes and the checked chain does not (`i16`, `i32`, and every floating point chain), the gap is the SIMD width times the check cost, up to 182x for the `f64` addition chain. +Where the baseline cannot vectorize either (the 64-bit and 128-bit chains, and the 8-bit chains whose narrow division defeats the vectorizer), the throwing types cost 1.0x to 1.9x, in line with the host benchmarks. +Where the baseline vectorizes and the throwing chain does not (`i16`, `i32`, and every floating point chain), the gap is the SIMD width times the check cost, up to roughly 235x for the `f64` chains. + +The deferred family recovers that gap almost entirely for floats: the branch-free `overflowing_*` operations vectorize with the baseline, so the same chains land at 1.1x to 2.9x, two orders of magnitude better than the throwing operators, while still detecting every error at the boundary (see xref:floats.adoc#floats_overflowing_arithmetic[Overflowing Arithmetic] and the `deferred_errors` accumulator in xref:sycl.adoc[]). +The integer `saturating_*` and `overflowing_*` chains improve `i16`/`i32` from about 9x to about 6x; the remaining gap is the divide by zero check that division and modulo still raise eagerly, which keeps a reporting branch in the kernel body. NOTE: The device numbers come from one machine and a single run each, so treat small differences as noise. On a discrete GPU a SYCL build should behave like the CUDA column; this machine only exposes the CPU device. diff --git a/doc/modules/ROOT/pages/cuda.adoc b/doc/modules/ROOT/pages/cuda.adoc index 1221f37..183b7d1 100644 --- a/doc/modules/ROOT/pages/cuda.adoc +++ b/doc/modules/ROOT/pages/cuda.adoc @@ -345,6 +345,11 @@ Be aware that other threads in the kernel may continue executing with incorrect * xref:examples.adoc#examples_cuda_error_handling[CUDA Error Handling] — shows how to use `device_error_context` to catch device-side overflow on the host and recover gracefully. * xref:examples.adoc#examples_cuda_error_handling_without_error_context[CUDA Error Handling Without Error Context] — demonstrates what happens when an overflow occurs on the device *without* `device_error_context`: the CUDA context is irrecoverably corrupted and no further kernels can be launched. +== Deferred Error Reporting + +The `deferred_errors` accumulator documented in xref:sycl.adoc#sycl_deferred_errors[the SYCL chapter] works identically under CUDA: accumulate the flags of the `overflowing_*` operations branch-free inside the kernel and call `report()` once per thread, and `device_error_context::synchronize()` surfaces the error on the host as usual. +On a GPU the throwing operators are already cheap (branches are well predicted), so the deferred mode matters mainly for `f32` heavy compute-bound kernels, where it removes the per-operation result classification. + == Benchmarks -xref:benchmarks.adoc#device_benchmarks[Dedicated device benchmarks] measure the overhead of the safety checks inside CUDA kernels, in both a memory-bound streaming shape and a compute-bound register chain. +xref:benchmarks.adoc#device_benchmarks[Dedicated device benchmarks] measure the overhead of the safety checks inside CUDA kernels, in both a memory-bound streaming shape and a compute-bound register chain, including the non-throwing `saturating_*` and `overflowing_*` forms. diff --git a/doc/modules/ROOT/pages/floats.adoc b/doc/modules/ROOT/pages/floats.adoc index cafc510..aed110e 100644 --- a/doc/modules/ROOT/pages/floats.adoc +++ b/doc/modules/ROOT/pages/floats.adoc @@ -187,6 +187,53 @@ In every operation, an operand that is a quiet or signaling NaN causes the opera NOTE: There is no `operator%` for `f32`/`f64`. C++ has no built-in floating-point remainder operator; use the `xref:cmath.adoc[fmod]` function instead. +[#floats_overflowing_arithmetic] +=== Overflowing Arithmetic + +[source,c++] +---- +// Overflowing arithmetic (raw IEEE result and error flag) +template +constexpr auto overflowing_add(float_basis lhs, float_basis rhs) noexcept + -> std::pair, bool>; + +template +constexpr auto overflowing_sub(float_basis lhs, float_basis rhs) noexcept + -> std::pair, bool>; + +template +constexpr auto overflowing_mul(float_basis lhs, float_basis rhs) noexcept + -> std::pair, bool>; + +template +constexpr auto overflowing_div(float_basis lhs, float_basis rhs) noexcept + -> std::pair, bool>; +---- + +The overflowing family is the non-throwing counterpart of the arithmetic operators, mirroring the integer `overflowing_*` functions. +The returned value is the raw IEEE result and the flag is `true` exactly when the corresponding operator would have thrown, which for every floating point operation means the result is non-finite. +The distinction between the exception types is not preserved: the flag only records that an error occurred. + +Because the functions are `noexcept` and branch-free, loops over them remain auto-vectorizable, which the throwing operators inherently prevent (a potential throw is an early exit the optimizer must preserve per element). +Accumulate the flags in an `unsigned` value and test once at a boundary of your choosing: + +[source,c++] +---- +auto flags {0U}; +for (std::size_t i {}; i < n; ++i) +{ + const auto r {overflowing_add(out[i], a[i])}; + out[i] = r.first; + flags |= static_cast(r.second); +} +if (flags != 0U) +{ + // At least one element produced a non-finite result +} +---- + +NOTE: Accumulate the flag in an `unsigned` value rather than a `bool`. Compilers recognize the integer OR reduction reliably; a `bool` accumulator can prevent the loop from vectorizing. + == Exception Behavior The following table summarizes the exceptional conditions and the exception each one produces. diff --git a/doc/modules/ROOT/pages/sycl.adoc b/doc/modules/ROOT/pages/sycl.adoc index ff3624a..93bfb4f 100644 --- a/doc/modules/ROOT/pages/sycl.adoc +++ b/doc/modules/ROOT/pages/sycl.adoc @@ -237,6 +237,48 @@ The error state is cleared *before* throwing, so after catching the exception th * xref:examples.adoc#examples_sycl[SYCL Device Support]: runs safe_numbers arithmetic on a SYCL device and verifies the results against the host. * xref:examples.adoc#examples_sycl_error_handling[SYCL Error Handling]: shows how to use `device_error_context` to catch a device-side error on the host. +[#sycl_deferred_errors] +== Deferred Error Reporting + +A CPU SYCL device executes kernels by vectorizing them across work items, and the per-operation reporting branch of the throwing operators inhibits exactly that transformation: the device benchmarks measure the branch-free baseline running up to two orders of magnitude faster than the checked kernel in compute-bound float loops. + +For hot kernels the library therefore offers a deferred mode built on the `overflowing_*` families (see xref:floats.adoc#floats_overflowing_arithmetic[Overflowing Arithmetic] for floats; the integer functions are documented with each integer type). +The `deferred_errors` accumulator collects the per-operation error flags without branching and raises them once per work item, through the same `device_global` machinery the operators use, so `device_error_context::synchronize()` surfaces the error on the host exactly as usual: + +[source, c++] +---- +#include +#include +#include + +using boost::safe_numbers::u32; + +void run(sycl::queue& q, const u32* in, u32* out, const int n) +{ + boost::safe_numbers::device_error_context ctx {q}; + + q.parallel_for(sycl::range<1>(static_cast(n)), [=](sycl::id<1> idx) + { + const auto i {idx[0]}; + + boost::safe_numbers::deferred_errors errs; + + // unwrap() accumulates the flag of each overflowing result branch-free + auto x {errs.unwrap(boost::safe_numbers::overflowing_add(in[i], in[i]))}; + x = errs.unwrap(boost::safe_numbers::overflowing_mul(x, in[i])); + + // The one branch per work item; records into the device error state + errs.report(__FILE__, __LINE__); + out[i] = x; + }); + + ctx.synchronize(); // rethrows on the host if any work item reported +} +---- + +The trade against the throwing operators is precision of the report: the error surfaces at the report() call rather than at the exact operation, and the individual categories are not preserved (a deferred error is always raised as overflow). +In exchange the kernel body stays vectorizable, which on CPU devices is the difference between roughly 2x and roughly 200x for floating point chains (see the benchmarks below). + == Benchmarks -xref:benchmarks.adoc#device_benchmarks[Dedicated device benchmarks] measure the overhead of the safety checks inside SYCL kernels, in both a memory-bound streaming shape and a compute-bound register chain. +xref:benchmarks.adoc#device_benchmarks[Dedicated device benchmarks] measure the overhead of the safety checks inside SYCL kernels, in both a memory-bound streaming shape and a compute-bound register chain, including the non-throwing `saturating_*` and `overflowing_*` forms. diff --git a/include/boost/safe_numbers/detail/float_basis.hpp b/include/boost/safe_numbers/detail/float_basis.hpp index 0c9225d..b5f8e90 100644 --- a/include/boost/safe_numbers/detail/float_basis.hpp +++ b/include/boost/safe_numbers/detail/float_basis.hpp @@ -1740,6 +1740,77 @@ BOOST_SAFE_NUMBERS_DEFINE_MIXED_FLOAT_OP("division", operator/) } // namespace boost::safe_numbers::detail +// ------------------------------ +// Deferred (overflowing) Math +// ------------------------------ + +namespace boost::safe_numbers { + +namespace detail::impl { + +// True exactly when the checked float operators would report res as an error: +// every one of them errors precisely on a non-finite result (the classifiers +// above only choose which exception to raise). The comparison is false for +// both inf and NAN, so the negation flags both without a branch. +template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto nonfinite_result(const T res) noexcept -> bool +{ + return !(constexpr_abs(res) <= (std::numeric_limits::max)()); +} + +} // namespace detail::impl + +// Non-throwing counterparts of the checked float operators for hot loops. The +// value is the raw IEEE result and the bool is true exactly when the checked +// operator would have thrown. The branch-free form keeps loops vectorizable: +// accumulate the flags in an unsigned value and test once at a boundary. + +template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto overflowing_add(const detail::float_basis lhs, + const detail::float_basis rhs) noexcept + -> std::pair, bool> +{ + const auto res {static_cast(static_cast(lhs) + static_cast(rhs))}; + return std::make_pair(detail::float_basis{res}, detail::impl::nonfinite_result(res)); +} + +BOOST_SAFE_NUMBERS_DEFINE_MIXED_FLOAT_OP("overflowing addition", overflowing_add) + +template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto overflowing_sub(const detail::float_basis lhs, + const detail::float_basis rhs) noexcept + -> std::pair, bool> +{ + const auto res {static_cast(static_cast(lhs) - static_cast(rhs))}; + return std::make_pair(detail::float_basis{res}, detail::impl::nonfinite_result(res)); +} + +BOOST_SAFE_NUMBERS_DEFINE_MIXED_FLOAT_OP("overflowing subtraction", overflowing_sub) + +template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto overflowing_mul(const detail::float_basis lhs, + const detail::float_basis rhs) noexcept + -> std::pair, bool> +{ + const auto res {static_cast(static_cast(lhs) * static_cast(rhs))}; + return std::make_pair(detail::float_basis{res}, detail::impl::nonfinite_result(res)); +} + +BOOST_SAFE_NUMBERS_DEFINE_MIXED_FLOAT_OP("overflowing multiplication", overflowing_mul) + +template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto overflowing_div(const detail::float_basis lhs, + const detail::float_basis rhs) noexcept + -> std::pair, bool> +{ + const auto res {static_cast(static_cast(lhs) / static_cast(rhs))}; + return std::make_pair(detail::float_basis{res}, detail::impl::nonfinite_result(res)); +} + +BOOST_SAFE_NUMBERS_DEFINE_MIXED_FLOAT_OP("overflowing division", overflowing_div) + +} // namespace boost::safe_numbers + #undef BOOST_SAFE_NUMBERS_DEFINE_MIXED_FLOAT_OP #endif // BOOST_SAFE_NUMBERS_DETAIL_FLOAT_BASIS_HPP diff --git a/include/boost/safe_numbers/detail/signed_integer_basis.hpp b/include/boost/safe_numbers/detail/signed_integer_basis.hpp index 3420994..8d2c81e 100644 --- a/include/boost/safe_numbers/detail/signed_integer_basis.hpp +++ b/include/boost/safe_numbers/detail/signed_integer_basis.hpp @@ -661,23 +661,9 @@ struct signed_add_helper const auto rhs_basis {static_cast(rhs)}; BasisType result {}; - #if BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_add_overflow) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X64_INTRIN) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN) - - if constexpr (!std::is_same_v) - { - #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) - - if (!std::is_constant_evaluated()) - { - const auto status {impl::signed_intrin_add(lhs_basis, rhs_basis, result)}; - return std::make_pair(result_type{result}, status != signed_overflow_status::no_error); - } - - #endif - } - - #endif - + // No overflow intrinsic here on purpose: the *.with.overflow intrinsics + // block loop vectorization when the flag is consumed as a value, while + // this compare form vectorizes and emits the same scalar code. const auto status {impl::signed_no_intrin_add(lhs_basis, rhs_basis, result)}; return std::make_pair(result_type{result}, status != signed_overflow_status::no_error); } @@ -698,25 +684,9 @@ struct signed_add_helper const auto rhs_basis {static_cast(rhs)}; BasisType result {}; - #if BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_add_overflow) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X64_INTRIN) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN) - - if constexpr (!std::is_same_v) - { - #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) - - if (!std::is_constant_evaluated()) - { - const auto status {impl::signed_intrin_add(lhs_basis, rhs_basis, result)}; - return status != signed_overflow_status::no_error - ? std::nullopt - : std::make_optional(result_type{result}); - } - - #endif - } - - #endif - + // No overflow intrinsic here on purpose: the *.with.overflow intrinsics + // block loop vectorization when the flag is consumed as a value, while + // this compare form vectorizes and emits the same scalar code. const auto status {impl::signed_no_intrin_add(lhs_basis, rhs_basis, result)}; return status != signed_overflow_status::no_error ? std::nullopt @@ -1293,23 +1263,9 @@ struct signed_sub_helper const auto rhs_basis {static_cast(rhs)}; BasisType result {}; - #if BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_sub_overflow) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X64_INTRIN) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN) - - if constexpr (!std::is_same_v) - { - #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) - - if (!std::is_constant_evaluated()) - { - const auto status {impl::signed_intrin_sub(lhs_basis, rhs_basis, result)}; - return std::make_pair(result_type{result}, status != signed_overflow_status::no_error); - } - - #endif - } - - #endif - + // No overflow intrinsic here on purpose: the *.with.overflow intrinsics + // block loop vectorization when the flag is consumed as a value, while + // this compare form vectorizes and emits the same scalar code. const auto status {impl::signed_no_intrin_sub(lhs_basis, rhs_basis, result)}; return std::make_pair(result_type{result}, status != signed_overflow_status::no_error); } @@ -1330,25 +1286,9 @@ struct signed_sub_helper const auto rhs_basis {static_cast(rhs)}; BasisType result {}; - #if BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_sub_overflow) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X64_INTRIN) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN) - - if constexpr (!std::is_same_v) - { - #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) - - if (!std::is_constant_evaluated()) - { - const auto status {impl::signed_intrin_sub(lhs_basis, rhs_basis, result)}; - return status != signed_overflow_status::no_error - ? std::nullopt - : std::make_optional(result_type{result}); - } - - #endif - } - - #endif - + // No overflow intrinsic here on purpose: the *.with.overflow intrinsics + // block loop vectorization when the flag is consumed as a value, while + // this compare form vectorizes and emits the same scalar code. const auto status {impl::signed_no_intrin_sub(lhs_basis, rhs_basis, result)}; return status != signed_overflow_status::no_error ? std::nullopt @@ -1897,27 +1837,9 @@ struct signed_mul_helper const auto rhs_basis {static_cast(rhs)}; BasisType result {}; - #if BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_mul_overflow) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X64_INTRIN) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN) - - // Route int128 through the intrin path only where the signed 128-bit - // fast path links (GCC, clang >= 14); elsewhere it uses no_intrin. - #if !defined(BOOST_SAFE_NUMBERS_HAS_INT128_SIGNED_INTRIN_MUL) - if constexpr (!std::is_same_v) - #endif - { - #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) - - if (!std::is_constant_evaluated()) - { - const auto status {impl::signed_intrin_mul(lhs_basis, rhs_basis, result)}; - return std::make_pair(result_type{result}, status != signed_overflow_status::no_error); - } - - #endif - } - - #endif - + // No overflow intrinsic here on purpose: the *.with.overflow intrinsics + // block loop vectorization when the flag is consumed as a value, while + // this compare form vectorizes and emits the same scalar code. const auto status {impl::signed_no_intrin_mul(lhs_basis, rhs_basis, result)}; return std::make_pair(result_type{result}, status != signed_overflow_status::no_error); } @@ -1938,29 +1860,9 @@ struct signed_mul_helper const auto rhs_basis {static_cast(rhs)}; BasisType result {}; - #if BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_mul_overflow) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X64_INTRIN) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN) - - // Route int128 through the intrin path only where the signed 128-bit - // fast path links (GCC, clang >= 14); elsewhere it uses no_intrin. - #if !defined(BOOST_SAFE_NUMBERS_HAS_INT128_SIGNED_INTRIN_MUL) - if constexpr (!std::is_same_v) - #endif - { - #if !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) - - if (!std::is_constant_evaluated()) - { - const auto status {impl::signed_intrin_mul(lhs_basis, rhs_basis, result)}; - return status != signed_overflow_status::no_error - ? std::nullopt - : std::make_optional(result_type{result}); - } - - #endif - } - - #endif - + // No overflow intrinsic here on purpose: the *.with.overflow intrinsics + // block loop vectorization when the flag is consumed as a value, while + // this compare form vectorizes and emits the same scalar code. const auto status {impl::signed_no_intrin_mul(lhs_basis, rhs_basis, result)}; return status != signed_overflow_status::no_error ? std::nullopt @@ -2773,6 +2675,7 @@ constexpr auto signed_integer_basis::operator--(int) namespace boost::safe_numbers { template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto saturating_add(const detail::signed_integer_basis lhs, const detail::signed_integer_basis rhs) noexcept -> detail::signed_integer_basis @@ -2783,6 +2686,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_SIGNED_INTEGER_OP("saturating addition", saturating_add) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto saturating_sub(const detail::signed_integer_basis lhs, const detail::signed_integer_basis rhs) noexcept -> detail::signed_integer_basis @@ -2793,6 +2697,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_SIGNED_INTEGER_OP("saturating subtraction", saturating_sub) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto saturating_mul(const detail::signed_integer_basis lhs, const detail::signed_integer_basis rhs) noexcept -> detail::signed_integer_basis @@ -2803,6 +2708,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_SIGNED_INTEGER_OP("saturating multiplication", saturating_mul) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto saturating_div(const detail::signed_integer_basis lhs, const detail::signed_integer_basis rhs) -> detail::signed_integer_basis @@ -2813,6 +2719,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_SIGNED_INTEGER_OP("saturating division", saturating_div) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto saturating_mod(const detail::signed_integer_basis lhs, const detail::signed_integer_basis rhs) -> detail::signed_integer_basis @@ -2827,6 +2734,7 @@ BOOST_SAFE_NUMBERS_DEFINE_MIXED_SIGNED_INTEGER_OP("saturating modulo", saturatin // ------------------------------ template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto overflowing_add(const detail::signed_integer_basis lhs, const detail::signed_integer_basis rhs) noexcept -> std::pair, bool> @@ -2837,6 +2745,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_SIGNED_INTEGER_OP("overflowing addition", overflowing_add) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto overflowing_sub(const detail::signed_integer_basis lhs, const detail::signed_integer_basis rhs) noexcept -> std::pair, bool> @@ -2847,6 +2756,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_SIGNED_INTEGER_OP("overflowing subtraction", overflowing_sub) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto overflowing_mul(const detail::signed_integer_basis lhs, const detail::signed_integer_basis rhs) noexcept -> std::pair, bool> @@ -2857,6 +2767,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_SIGNED_INTEGER_OP("overflowing multiplication", overflowing_mul) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto overflowing_div(const detail::signed_integer_basis lhs, const detail::signed_integer_basis rhs) -> std::pair, bool> @@ -2867,6 +2778,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_SIGNED_INTEGER_OP("overflowing division", overflowing_div) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto overflowing_mod(const detail::signed_integer_basis lhs, const detail::signed_integer_basis rhs) -> std::pair, bool> @@ -2881,6 +2793,7 @@ BOOST_SAFE_NUMBERS_DEFINE_MIXED_SIGNED_INTEGER_OP("overflowing modulo", overflow // ------------------------------ template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto checked_add(const detail::signed_integer_basis lhs, const detail::signed_integer_basis rhs) noexcept -> std::optional> @@ -2891,6 +2804,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_SIGNED_INTEGER_OP("checked addition", checked_add) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto checked_sub(const detail::signed_integer_basis lhs, const detail::signed_integer_basis rhs) noexcept -> std::optional> @@ -2901,6 +2815,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_SIGNED_INTEGER_OP("checked subtraction", checked_sub) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto checked_mul(const detail::signed_integer_basis lhs, const detail::signed_integer_basis rhs) noexcept -> std::optional> @@ -2911,6 +2826,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_SIGNED_INTEGER_OP("checked multiplication", checked_mul) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto checked_div(const detail::signed_integer_basis lhs, const detail::signed_integer_basis rhs) noexcept -> std::optional> @@ -2921,6 +2837,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_SIGNED_INTEGER_OP("checked division", checked_div) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto checked_mod(const detail::signed_integer_basis lhs, const detail::signed_integer_basis rhs) noexcept -> std::optional> diff --git a/include/boost/safe_numbers/detail/unsigned_integer_basis.hpp b/include/boost/safe_numbers/detail/unsigned_integer_basis.hpp index cd8e1c1..807edff 100644 --- a/include/boost/safe_numbers/detail/unsigned_integer_basis.hpp +++ b/include/boost/safe_numbers/detail/unsigned_integer_basis.hpp @@ -642,19 +642,9 @@ struct add_helper const auto rhs_basis {static_cast(rhs)}; BasisType res {}; - if constexpr (!std::is_same_v) - { - #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_add_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_addcarry_u64) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN)) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) - - if (!std::is_constant_evaluated()) - { - const auto overflowed {impl::unsigned_intrin_add(lhs_basis, rhs_basis, res)}; - return std::make_pair(result_type{res}, overflowed); - } - - #endif - } - + // No overflow intrinsic here on purpose: the *.with.overflow intrinsics + // block loop vectorization when the flag is consumed as a value, while + // this compare form vectorizes and emits the same scalar code. const auto overflowed {impl::unsigned_no_intrin_add(lhs_basis, rhs_basis, res)}; return std::make_pair(result_type{res}, overflowed); } @@ -675,19 +665,9 @@ struct add_helper const auto rhs_basis {static_cast(rhs)}; BasisType res {}; - if constexpr (!std::is_same_v) - { - #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_add_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_addcarry_u64) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN)) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) - - if (!std::is_constant_evaluated()) - { - const auto overflowed {impl::unsigned_intrin_add(lhs_basis, rhs_basis, res)}; - return overflowed ? std::nullopt : std::make_optional(result_type{res}); - } - - #endif - } - + // No overflow intrinsic here on purpose: the *.with.overflow intrinsics + // block loop vectorization when the flag is consumed as a value, while + // this compare form vectorizes and emits the same scalar code. const auto overflowed {impl::unsigned_no_intrin_add(lhs_basis, rhs_basis, res)}; return overflowed ? std::nullopt : std::make_optional(result_type{res}); } @@ -1167,19 +1147,9 @@ struct sub_helper const auto rhs_basis {static_cast(rhs)}; BasisType res {}; - if constexpr (!std::is_same_v) - { - #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_sub_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_subborrow_u64) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN)) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) - - if (!std::is_constant_evaluated()) - { - const auto underflowed {impl::unsigned_intrin_sub(lhs_basis, rhs_basis, res)}; - return std::make_pair(result_type{res}, underflowed); - } - - #endif - } - + // No overflow intrinsic here on purpose: the *.with.overflow intrinsics + // block loop vectorization when the flag is consumed as a value, while + // this compare form vectorizes and emits the same scalar code. const auto underflowed {impl::unsigned_no_intrin_sub(lhs_basis, rhs_basis, res)}; return std::make_pair(result_type{res}, underflowed); } @@ -1200,19 +1170,9 @@ struct sub_helper const auto rhs_basis {static_cast(rhs)}; BasisType res {}; - if constexpr (!std::is_same_v) - { - #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_sub_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_subborrow_u64) || defined(BOOST_SAFENUMBERS_HAS_WINDOWS_X86_INTRIN)) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) - - if (!std::is_constant_evaluated()) - { - const auto underflowed {impl::unsigned_intrin_sub(lhs_basis, rhs_basis, res)}; - return underflowed ? std::nullopt : std::make_optional(result_type{res}); - } - - #endif - } - + // No overflow intrinsic here on purpose: the *.with.overflow intrinsics + // block loop vectorization when the flag is consumed as a value, while + // this compare form vectorizes and emits the same scalar code. const auto underflowed {impl::unsigned_no_intrin_sub(lhs_basis, rhs_basis, res)}; return underflowed ? std::nullopt : std::make_optional(result_type{res}); } @@ -1509,22 +1469,9 @@ struct mul_helper const auto rhs_basis {static_cast(rhs)}; BasisType res {}; - // We have an intrin path, but only with __builtin_mul_overflow - #if BOOST_SAFE_NUMBERS_HAS_BUILTIN(_umul128) || !defined(BOOST_SAFE_NUMBERS_DETAIL_INT128_HAS_INT128) - if constexpr (!std::is_same_v) - #endif - { - #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_mul_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_umul128)) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) - - if (!std::is_constant_evaluated()) - { - const auto overflowed {impl::unsigned_intrin_mul(lhs_basis, rhs_basis, res)}; - return std::make_pair(result_type{res}, overflowed); - } - - #endif - } - + // No overflow intrinsic here on purpose: the *.with.overflow intrinsics + // block loop vectorization when the flag is consumed as a value, while + // this compare form vectorizes and emits the same scalar code. const auto overflowed {impl::no_intrin_mul(lhs_basis, rhs_basis, res)}; return std::make_pair(result_type{res}, overflowed); } @@ -1545,22 +1492,9 @@ struct mul_helper const auto rhs_basis {static_cast(rhs)}; BasisType res {}; - // We have an intrin path, but only with __builtin_mul_overflow - #if BOOST_SAFE_NUMBERS_HAS_BUILTIN(_umul128) || !defined(BOOST_SAFE_NUMBERS_DETAIL_INT128_HAS_INT128) - if constexpr (!std::is_same_v) - #endif - { - #if (BOOST_SAFE_NUMBERS_HAS_BUILTIN(__builtin_mul_overflow) || BOOST_SAFE_NUMBERS_HAS_BUILTIN(_umul128)) && !defined(BOOST_SAFE_NUMBERS_HAS_GPU_SUPPORT) - - if (!std::is_constant_evaluated()) - { - const auto overflowed {impl::unsigned_intrin_mul(lhs_basis, rhs_basis, res)}; - return overflowed ? std::nullopt : std::make_optional(result_type{res}); - } - - #endif - } - + // No overflow intrinsic here on purpose: the *.with.overflow intrinsics + // block loop vectorization when the flag is consumed as a value, while + // this compare form vectorizes and emits the same scalar code. const auto overflowed {impl::no_intrin_mul(lhs_basis, rhs_basis, res)}; return overflowed ? std::nullopt : std::make_optional(result_type{res}); } @@ -2055,6 +1989,7 @@ constexpr auto unsigned_integer_basis::operator--(int) template struct shl_helper { + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] static constexpr auto apply(const unsigned_integer_basis lhs, const unsigned_integer_basis rhs) noexcept(Policy != overflow_policy::throw_exception) @@ -2105,6 +2040,7 @@ struct shl_helper template struct shl_helper { + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] static constexpr auto apply(const unsigned_integer_basis lhs, const unsigned_integer_basis rhs) noexcept -> std::pair, bool> @@ -2131,6 +2067,7 @@ struct shl_helper template struct shl_helper { + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] static constexpr auto apply(const unsigned_integer_basis lhs, const unsigned_integer_basis rhs) noexcept -> std::optional> @@ -2148,6 +2085,7 @@ struct shl_helper }; template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto shl_impl(const unsigned_integer_basis lhs, const unsigned_integer_basis rhs) noexcept(Policy == overflow_policy::saturate || Policy == overflow_policy::overflow_tuple || Policy == overflow_policy::checked || Policy == overflow_policy::strict) @@ -2163,6 +2101,7 @@ template template struct shr_helper { + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] static constexpr auto apply(const unsigned_integer_basis lhs, const unsigned_integer_basis rhs) noexcept(Policy != overflow_policy::throw_exception) @@ -2211,6 +2150,7 @@ struct shr_helper template struct shr_helper { + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] static constexpr auto apply(const unsigned_integer_basis lhs, const unsigned_integer_basis rhs) noexcept -> std::pair, bool> @@ -2234,6 +2174,7 @@ struct shr_helper template struct shr_helper { + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] static constexpr auto apply(const unsigned_integer_basis lhs, const unsigned_integer_basis rhs) noexcept -> std::optional> @@ -2249,6 +2190,7 @@ struct shr_helper }; template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto shr_impl(const unsigned_integer_basis lhs, const unsigned_integer_basis rhs) noexcept(Policy == overflow_policy::saturate || Policy == overflow_policy::overflow_tuple || Policy == overflow_policy::checked || Policy == overflow_policy::strict) @@ -2265,6 +2207,7 @@ template namespace boost::safe_numbers { template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto saturating_add(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) noexcept -> detail::unsigned_integer_basis @@ -2275,6 +2218,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_UNSIGNED_INTEGER_OP("saturating addition", saturating_add) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto saturating_sub(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) noexcept -> detail::unsigned_integer_basis @@ -2285,6 +2229,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_UNSIGNED_INTEGER_OP("saturating subtraction", saturating_sub) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto saturating_mul(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) noexcept -> detail::unsigned_integer_basis @@ -2295,6 +2240,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_UNSIGNED_INTEGER_OP("saturating multiplication", saturating_mul) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto saturating_div(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) -> detail::unsigned_integer_basis @@ -2305,6 +2251,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_UNSIGNED_INTEGER_OP("saturating division", saturating_div) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto saturating_mod(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) -> detail::unsigned_integer_basis @@ -2315,6 +2262,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_UNSIGNED_INTEGER_OP("saturating modulo", saturating_mod) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto overflowing_add(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) noexcept -> std::pair, bool> @@ -2325,6 +2273,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_UNSIGNED_INTEGER_OP("overflowing addition", overflowing_add) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto overflowing_sub(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) noexcept -> std::pair, bool> @@ -2335,6 +2284,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_UNSIGNED_INTEGER_OP("overflowing subtraction", overflowing_sub) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto overflowing_mul(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) noexcept -> std::pair, bool> @@ -2345,6 +2295,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_UNSIGNED_INTEGER_OP("overflowing multiplication", overflowing_mul) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto overflowing_div(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) -> std::pair, bool> @@ -2355,6 +2306,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_UNSIGNED_INTEGER_OP("overflowing division", overflowing_div) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto overflowing_mod(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) -> std::pair, bool> @@ -2365,6 +2317,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_UNSIGNED_INTEGER_OP("overflowing modulo", overflowing_mod) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto checked_add(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) noexcept -> std::optional> @@ -2375,6 +2328,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_UNSIGNED_INTEGER_OP("checked addition", checked_add) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto checked_sub(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) noexcept -> std::optional> @@ -2385,6 +2339,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_UNSIGNED_INTEGER_OP("checked subtraction", checked_sub) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto checked_mul(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) noexcept -> std::optional> @@ -2395,6 +2350,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_UNSIGNED_INTEGER_OP("checked multiplication", checked_mul) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto checked_div(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) noexcept -> std::optional> @@ -2405,6 +2361,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_UNSIGNED_INTEGER_OP("checked division", checked_div) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto checked_mod(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) noexcept -> std::optional> @@ -2487,6 +2444,7 @@ BOOST_SAFE_NUMBERS_DEFINE_MIXED_UNSIGNED_INTEGER_OP("widening mul", widening_mul // ------------------------------ template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto saturating_shl(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) noexcept -> detail::unsigned_integer_basis @@ -2497,6 +2455,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_UNSIGNED_INTEGER_OP("saturating left shift", saturating_shl) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto saturating_shr(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) noexcept -> detail::unsigned_integer_basis @@ -2511,6 +2470,7 @@ BOOST_SAFE_NUMBERS_DEFINE_MIXED_UNSIGNED_INTEGER_OP("saturating right shift", sa // ------------------------------ template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto overflowing_shl(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) noexcept -> std::pair, bool> @@ -2521,6 +2481,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_UNSIGNED_INTEGER_OP("overflowing left shift", overflowing_shl) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto overflowing_shr(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) noexcept -> std::pair, bool> @@ -2535,6 +2496,7 @@ BOOST_SAFE_NUMBERS_DEFINE_MIXED_UNSIGNED_INTEGER_OP("overflowing right shift", o // ------------------------------ template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto checked_shl(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) noexcept -> std::optional> @@ -2545,6 +2507,7 @@ template BOOST_SAFE_NUMBERS_DEFINE_MIXED_UNSIGNED_INTEGER_OP("checked left shift", checked_shl) template +BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto checked_shr(const detail::unsigned_integer_basis lhs, const detail::unsigned_integer_basis rhs) noexcept -> std::optional> diff --git a/include/boost/safe_numbers/device_error_reporting.hpp b/include/boost/safe_numbers/device_error_reporting.hpp index e6e0011..4efab42 100644 --- a/include/boost/safe_numbers/device_error_reporting.hpp +++ b/include/boost/safe_numbers/device_error_reporting.hpp @@ -25,4 +25,60 @@ # include #endif +#if (defined(__CUDACC__) && defined(BOOST_SAFE_NUMBERS_ENABLE_CUDA)) || defined(BOOST_SAFE_NUMBERS_ENABLE_SYCL) + +#ifndef BOOST_SAFE_NUMBERS_BUILD_MODULE +#include +#endif + +namespace boost::safe_numbers { + +// Deferred device error reporting for kernels built on the overflowing_* +// families. The per-operation flags accumulate branch-free, so the kernel +// body stays vectorizable on CPU style devices where the per-operation report +// branch of the throwing operators inhibits SPMD vectorization. report() is +// the one branch per work item: on the device it records into the same error +// state that device_error_context::synchronize() rethrows from, and on the +// host it throws immediately. The individual error categories are not +// preserved; a deferred error always surfaces as overflow. +class deferred_errors +{ +public: + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr void accumulate(const bool had_error) noexcept + { + flags_ |= static_cast(had_error); + } + + // Accumulates the flag of an overflowing_* result and returns its value + template + BOOST_SAFE_NUMBERS_HOST_DEVICE constexpr auto unwrap(const std::pair result) noexcept -> T + { + accumulate(result.second); + return result.first; + } + + BOOST_SAFE_NUMBERS_HOST_DEVICE [[nodiscard]] constexpr auto any() const noexcept -> bool + { + return flags_ != 0U; + } + + // Call once at the end of the kernel body + BOOST_SAFE_NUMBERS_HOST_DEVICE void report(const char* file = "deferred_errors", + const int line = 0, + const char* expression = "deferred checked arithmetic") const + { + if (flags_ != 0U) + { + detail::report_device_error(detail::exception_type::overflow, file, line, expression); + } + } + +private: + unsigned flags_ {0U}; +}; + +} // namespace boost::safe_numbers + +#endif + #endif // BOOST_SAFE_NUMBERS_DEVICE_ERROR_REPORTING_HPP diff --git a/test/Jamfile b/test/Jamfile index b228dd1..43c7929 100644 --- a/test/Jamfile +++ b/test/Jamfile @@ -61,6 +61,7 @@ run test_float_addition.cpp ; run test_float_subtraction.cpp ; run test_float_mul.cpp ; run test_float_div.cpp ; +run test_float_overflowing.cpp ; run test_float_streaming.cpp ; run test_float_std_format.cpp ; run test_float_fmt_format.cpp ; diff --git a/test/benchmarks/benchmark_cuda_operations.cu b/test/benchmarks/benchmark_cuda_operations.cu index 20dfae6..1602e3c 100644 --- a/test/benchmarks/benchmark_cuda_operations.cu +++ b/test/benchmarks/benchmark_cuda_operations.cu @@ -301,6 +301,7 @@ int run_int_family(const char* name, boost::safe_numbers::device_error_context& failures += run_stream_op(ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg); failures += run_stream_op(ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg); failures += run_stream_op(ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg); + failures += run_stream_op(ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg); } { @@ -308,6 +309,8 @@ int run_int_family(const char* name, boost::safe_numbers::device_error_context& device_bench::print_table_header(name, "register chain of checked ops (compute-bound)"); failures += run_chain(ctx, in, cfg); + failures += run_chain(ctx, in, cfg); + failures += run_chain(ctx, in, cfg); } return failures; @@ -354,6 +357,11 @@ int run_real_family(const char* name, boost::safe_numbers::device_error_context& const auto div_in {device_bench::gen_real_chain(cfg.chain_threads, device_bench::real_chain_kind::divisor)}; failures += run_chain(ctx, div_in, cfg); + + failures += run_chain(ctx, add_in, cfg); + failures += run_chain(ctx, add_in, cfg); + failures += run_chain(ctx, mul_in, cfg); + failures += run_chain(ctx, div_in, cfg); } return failures; diff --git a/test/benchmarks/benchmark_sycl_operations.cpp b/test/benchmarks/benchmark_sycl_operations.cpp index d21a422..8f42d97 100644 --- a/test/benchmarks/benchmark_sycl_operations.cpp +++ b/test/benchmarks/benchmark_sycl_operations.cpp @@ -280,6 +280,7 @@ int run_int_family(const char* name, sycl::queue& q, boost::safe_numbers::device failures += run_stream_op(q, ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg); failures += run_stream_op(q, ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg); failures += run_stream_op(q, ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg); + failures += run_stream_op(q, ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg); } { @@ -287,6 +288,8 @@ int run_int_family(const char* name, sycl::queue& q, boost::safe_numbers::device device_bench::print_table_header(name, "register chain of checked ops (compute-bound)"); failures += run_chain(q, ctx, in, cfg); + failures += run_chain(q, ctx, in, cfg); + failures += run_chain(q, ctx, in, cfg); } return failures; @@ -333,6 +336,11 @@ int run_real_family(const char* name, sycl::queue& q, boost::safe_numbers::devic const auto div_in {device_bench::gen_real_chain(cfg.chain_threads, device_bench::real_chain_kind::divisor)}; failures += run_chain(q, ctx, div_in, cfg); + + failures += run_chain(q, ctx, add_in, cfg); + failures += run_chain(q, ctx, add_in, cfg); + failures += run_chain(q, ctx, mul_in, cfg); + failures += run_chain(q, ctx, div_in, cfg); } return failures; diff --git a/test/benchmarks/device_bench.hpp b/test/benchmarks/device_bench.hpp index 9a0f3fd..cde400f 100644 --- a/test/benchmarks/device_bench.hpp +++ b/test/benchmarks/device_bench.hpp @@ -120,6 +120,26 @@ struct op_mod } }; +// Saturating addition for the safe type against the raw addition baseline: +// shows what the noexcept value policy costs relative to unchecked arithmetic +struct op_sat_add +{ + static constexpr const char* name {"sat_add"}; + + template + BOOST_SAFE_NUMBERS_HOST_DEVICE static T apply(const T a, const T b) + { + if constexpr (requires { typename T::basis_type; }) + { + return boost::safe_numbers::saturating_add(a, b); + } + else + { + return static_cast(a + b); + } + } +}; + // --------------------------------------------------------------------------- // Compute-bound chains // @@ -165,6 +185,84 @@ struct int_mixed_chain } }; +// Saturating variant of the mixed chain. The safe type runs the noexcept +// saturating_* functions; the basis type runs the raw chain, so the ratio +// shows the cost of the value policy against unchecked arithmetic. The +// operand ranges never saturate, so the results match the raw chain exactly. +template +struct int_mixed_sat_chain +{ + static constexpr int ops_per_cycle {8}; + static constexpr const char* name {"mixed (sat)"}; + + BOOST_SAFE_NUMBERS_HOST_DEVICE static T cycle(T x, const T* y, const T aux) + { + if constexpr (requires { typename T::basis_type; }) + { + static_cast(aux); + namespace sn = boost::safe_numbers; + x = sn::saturating_mod(x, y[0]); + x = sn::saturating_add(x, y[1]); + x = sn::saturating_div(x, y[2]); + x = sn::saturating_sub(x, y[3]); + x = sn::saturating_mul(x, y[4]); + x = sn::saturating_mod(x, y[5]); + x = sn::saturating_add(x, y[6]); + x = sn::saturating_div(x, y[7]); + return x; + } + else + { + return int_mixed_chain::cycle(x, y, aux); + } + } +}; + +// Overflowing (deferred flag) variant of the mixed chain. Each step keeps its +// error flag alive by accumulating into an unsigned value that is folded back +// into the result once per cycle, mirroring how the family is meant to be +// used in kernels. The flags are always zero here, so the fold adds zero and +// the results match the raw chain exactly. The fold is not counted in +// ops_per_cycle, so the reported per-op cost is slightly conservative. +template +struct int_mixed_ovf_chain +{ + static constexpr int ops_per_cycle {8}; + static constexpr const char* name {"mixed (ovf)"}; + + BOOST_SAFE_NUMBERS_HOST_DEVICE static T cycle(T x, const T* y, const T aux) + { + if constexpr (requires { typename T::basis_type; }) + { + static_cast(aux); + namespace sn = boost::safe_numbers; + using basis = typename T::basis_type; + + auto flags {0U}; + const auto step {[&flags](const auto r) + { + flags |= static_cast(r.second); + return r.first; + }}; + + x = step(sn::overflowing_mod(x, y[0])); + x = step(sn::overflowing_add(x, y[1])); + x = step(sn::overflowing_div(x, y[2])); + x = step(sn::overflowing_sub(x, y[3])); + x = step(sn::overflowing_mul(x, y[4])); + x = step(sn::overflowing_mod(x, y[5])); + x = step(sn::overflowing_add(x, y[6])); + x = step(sn::overflowing_div(x, y[7])); + + return boost::safe_numbers::overflowing_add(x, T{static_cast(flags)}).first; + } + else + { + return int_mixed_chain::cycle(x, y, aux); + } + } +}; + // aux is the precomputed sum of the eight operands, so the restore step // returns x to its starting point (up to rounding, which stays bounded) template @@ -252,6 +350,58 @@ struct real_div_chain } }; +// Overflowing (deferred flag) variants of the float chains. The safe type +// runs the noexcept overflowing_* family with the flags accumulated and +// folded back into the value once per cycle; the basis type runs the raw +// chain. The values never leave the finite range, so the flags stay zero and +// the results match the raw chain exactly. +#define BOOST_SAFE_NUMBERS_BENCH_REAL_OVF_CHAIN(CHAIN_NAME, LABEL, STEP_FN, RESTORE_FN, FALLBACK) \ +template \ +struct CHAIN_NAME \ +{ \ + static constexpr int ops_per_cycle {9}; \ + static constexpr const char* name {LABEL}; \ + \ + BOOST_SAFE_NUMBERS_HOST_DEVICE static T cycle(T x, const T* y, const T aux) \ + { \ + if constexpr (requires { typename T::basis_type; }) \ + { \ + namespace sn = boost::safe_numbers; \ + using basis = typename T::basis_type; \ + \ + auto flags {0U}; \ + const auto step {[&flags](const auto r) \ + { \ + flags |= static_cast(r.second); \ + return r.first; \ + }}; \ + \ + x = step(sn::STEP_FN(x, y[0])); \ + x = step(sn::STEP_FN(x, y[1])); \ + x = step(sn::STEP_FN(x, y[2])); \ + x = step(sn::STEP_FN(x, y[3])); \ + x = step(sn::STEP_FN(x, y[4])); \ + x = step(sn::STEP_FN(x, y[5])); \ + x = step(sn::STEP_FN(x, y[6])); \ + x = step(sn::STEP_FN(x, y[7])); \ + x = step(sn::RESTORE_FN(x, aux)); \ + \ + return sn::overflowing_add(x, T{static_cast(flags)}).first; \ + } \ + else \ + { \ + return FALLBACK::cycle(x, y, aux); \ + } \ + } \ +}; + +BOOST_SAFE_NUMBERS_BENCH_REAL_OVF_CHAIN(real_ovf_add_chain, "add (ovf)", overflowing_add, overflowing_sub, real_add_chain) +BOOST_SAFE_NUMBERS_BENCH_REAL_OVF_CHAIN(real_ovf_sub_chain, "sub (ovf)", overflowing_sub, overflowing_add, real_sub_chain) +BOOST_SAFE_NUMBERS_BENCH_REAL_OVF_CHAIN(real_ovf_mul_chain, "mul (ovf)", overflowing_mul, overflowing_div, real_mul_chain) +BOOST_SAFE_NUMBERS_BENCH_REAL_OVF_CHAIN(real_ovf_div_chain, "div (ovf)", overflowing_div, overflowing_mul, real_div_chain) + +#undef BOOST_SAFE_NUMBERS_BENCH_REAL_OVF_CHAIN + // --------------------------------------------------------------------------- // Host-side input generation. Every generator fills a builtin vector and a // safe vector from the same random draws so both contenders see identical diff --git a/test/cuda_jamfile b/test/cuda_jamfile index 8990941..f3a5973 100644 --- a/test/cuda_jamfile +++ b/test/cuda_jamfile @@ -12,6 +12,7 @@ project : requirements # Test the error handler mechanisms run test_cuda_error_handling.cu ; run test_cuda_invalid_argument.cu ; +run test_cuda_deferred_errors.cu ; # u8 tests run test_cuda_u8_add.cu ; diff --git a/test/sycl_jamfile b/test/sycl_jamfile index c8838fd..e3496f1 100644 --- a/test/sycl_jamfile +++ b/test/sycl_jamfile @@ -58,6 +58,7 @@ run test_device_error_context_reuse_sycl.cpp ; run test_device_error_single_context_sycl.cpp ; run test_device_error_mapping_sycl.cpp ; run test_device_error_workitem_id_sycl.cpp ; +run test_deferred_errors_sycl.cpp ; # Free functions - bit manipulation run test_has_single_bit_sycl.cpp ; diff --git a/test/test_cuda_deferred_errors.cu b/test/test_cuda_deferred_errors.cu new file mode 100644 index 0000000..efb62aa --- /dev/null +++ b/test/test_cuda_deferred_errors.cu @@ -0,0 +1,103 @@ +// 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 deferred_errors: overflowing_* ops accumulate their flags +// branch-free and report() surfaces the result once per thread through the +// same machinery as the throwing operators. + +#include +#include +#include +#include +#include +#include "cuda_managed_ptr.hpp" + +#include + +using test_type = boost::safe_numbers::u32; +using basis_type = test_type::basis_type; + +__global__ void deferred_kernel(const test_type *in, test_type *out, int numElements) +{ + int i = blockDim.x * blockIdx.x + threadIdx.x; + + if (i < numElements) + { + boost::safe_numbers::deferred_errors errs; + + auto x {errs.unwrap(boost::safe_numbers::overflowing_add(in[i], in[i]))}; + x = errs.unwrap(boost::safe_numbers::overflowing_mul(x, test_type{basis_type{2}})); + x = errs.unwrap(boost::safe_numbers::overflowing_sub(x, in[i])); + + errs.report(__FILE__, __LINE__); + out[i] = x; + } +} + +int main(void) +{ + int numElements = 1024; + int threadsPerBlock = 256; + int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock; + + cuda_managed_ptr input_vector(numElements); + cuda_managed_ptr output_vector(numElements); + + boost::safe_numbers::device_error_context ctx; + + // Happy path: values stay in range, no flag is ever set + for (int i = 0; i < numElements; ++i) + { + input_vector[i] = test_type{static_cast(i % 1000 + 1)}; + } + + deferred_kernel<<>>(input_vector.get(), output_vector.get(), numElements); + + try + { + ctx.synchronize(); + } + catch (const std::exception& e) + { + std::cerr << "Unexpected error on the happy path: " << e.what() << std::endl; + return EXIT_FAILURE; + } + + for (int i = 0; i < numElements; ++i) + { + const auto v {static_cast(i % 1000 + 1)}; + if (static_cast(output_vector[i]) != static_cast(v * 4 - v)) + { + std::cerr << "Result verification failed at element " << i << "!" << std::endl; + return EXIT_FAILURE; + } + } + + // Error path: max + max sets the deferred flag and report() records it + for (int i = 0; i < numElements; ++i) + { + input_vector[i] = test_type{(std::numeric_limits::max)()}; + } + + deferred_kernel<<>>(input_vector.get(), output_vector.get(), numElements); + + try + { + ctx.synchronize(); + std::cerr << "Expected exception from deferred report but none was thrown!" << std::endl; + return EXIT_FAILURE; + } + catch (const std::overflow_error& e) + { + std::cerr << "Caught expected error: " << e.what() << std::endl; + std::cerr << "Test PASSED\n"; + return 0; + } + catch (const std::exception& e) + { + std::cerr << "Wrong exception type: " << e.what() << std::endl; + return EXIT_FAILURE; + } +} diff --git a/test/test_deferred_errors_sycl.cpp b/test/test_deferred_errors_sycl.cpp new file mode 100644 index 0000000..6977efe --- /dev/null +++ b/test/test_deferred_errors_sycl.cpp @@ -0,0 +1,123 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// Exercises deferred_errors on a SYCL device: overflowing_* ops accumulate +// their flags branch-free inside the kernel and report() surfaces the result +// once per work item through the same device_global machinery that the +// throwing operators use. + +#include "sycl_test.hpp" + +using test_type = boost::safe_numbers::u32; +using basis_type = test_type::basis_type; + +namespace { + +sycl::event launch_deferred(sycl::queue& q, const test_type* in, test_type* out, const int n) +{ + return q.parallel_for(sycl::range<1>(static_cast(n)), [=](sycl::id<1> idx) + { + const auto i {idx[0]}; + + boost::safe_numbers::deferred_errors errs; + + auto x {errs.unwrap(boost::safe_numbers::overflowing_add(in[i], in[i]))}; + x = errs.unwrap(boost::safe_numbers::overflowing_mul(x, test_type{basis_type{2}})); + x = errs.unwrap(boost::safe_numbers::overflowing_sub(x, in[i])); + + errs.report(__FILE__, __LINE__); + out[i] = x; + }); +} + +int test_happy_path() +{ + constexpr int n {1024}; + + sycl::queue q; + boost::safe_numbers::device_error_context ctx {q}; + + auto* in {sycl::malloc_shared(n, q)}; + auto* out {sycl::malloc_shared(n, q)}; + for (int i {0}; i < n; ++i) + { + in[i] = test_type{static_cast(i % 1000 + 1)}; + } + + launch_deferred(q, in, out, n); + + auto rc {EXIT_SUCCESS}; + try + { + ctx.synchronize(); + + auto failures {0}; + for (int i {0}; i < n; ++i) + { + const auto v {static_cast(i % 1000 + 1)}; + if (static_cast(out[i]) != static_cast(v * 4 - v)) + { + ++failures; + } + } + std::cout << "deferred happy path: " << (failures == 0 ? "PASSED" : "FAILED") << "\n"; + rc = failures == 0 ? EXIT_SUCCESS : EXIT_FAILURE; + } + catch (const std::exception& e) + { + std::cerr << "deferred happy path: FAILED - unexpected error: " << e.what() << "\n"; + rc = EXIT_FAILURE; + } + + sycl::free(in, q); + sycl::free(out, q); + return rc; +} + +int test_error_path() +{ + constexpr int n {1024}; + + sycl::queue q; + boost::safe_numbers::device_error_context ctx {q}; + + auto* in {sycl::malloc_shared(n, q)}; + auto* out {sycl::malloc_shared(n, q)}; + for (int i {0}; i < n; ++i) + { + in[i] = test_type{(std::numeric_limits::max)()}; + } + + launch_deferred(q, in, out, n); + + auto rc {EXIT_FAILURE}; + try + { + ctx.synchronize(); + std::cerr << "deferred error path: FAILED - no exception thrown\n"; + } + catch (const std::overflow_error& e) + { + std::cout << "deferred error path: PASSED (" << e.what() << ")\n"; + rc = EXIT_SUCCESS; + } + catch (const std::exception& e) + { + std::cerr << "deferred error path: FAILED - wrong exception type: " << e.what() << "\n"; + } + + sycl::free(in, q); + sycl::free(out, q); + return rc; +} + +} // namespace + +int main() +{ + auto rc {EXIT_SUCCESS}; + rc |= test_happy_path(); + rc |= test_error_path(); + return rc; +} diff --git a/test/test_float_overflowing.cpp b/test/test_float_overflowing.cpp new file mode 100644 index 0000000..3943c9f --- /dev/null +++ b/test/test_float_overflowing.cpp @@ -0,0 +1,119 @@ +// Copyright 2026 Matt Borland +// Distributed under the Boost Software License, Version 1.0. +// https://www.boost.org/LICENSE_1_0.txt +// +// The overflowing float family must agree with the throwing operators on +// every input: the bool is true exactly when the operator throws, and when it +// does not throw the value is bit-identical to the operator's result. + +#include +#include + +#if defined(__clang__) +# pragma clang diagnostic push +# pragma clang diagnostic ignored "-Wfloat-equal" +#elif defined(__GNUC__) +# pragma GCC diagnostic push +# pragma GCC diagnostic ignored "-Wfloat-equal" +#endif + +#include +#include +#include +#include +#include +#include + +using namespace boost::safe_numbers; + +template +using bits_type = std::conditional_t; + +template +void test_parity(OvfOp ovf, ThrowOp thr) +{ + using basis_type = typename SafeT::basis_type; + + const std::vector specials { + basis_type{0}, + static_cast(-0.0), + basis_type{1}, + static_cast(-1), + static_cast(0.5), + basis_type{100}, + (std::numeric_limits::max)(), + -(std::numeric_limits::max)(), + (std::numeric_limits::min)(), + std::numeric_limits::denorm_min(), + std::numeric_limits::infinity(), + -std::numeric_limits::infinity(), + std::numeric_limits::quiet_NaN(), + }; + + for (const auto a : specials) + { + for (const auto b : specials) + { + const auto r {ovf(SafeT{a}, SafeT{b})}; + + auto threw {false}; + auto expected {SafeT{basis_type{0}}}; + try + { + expected = thr(SafeT{a}, SafeT{b}); + } + catch (const std::exception&) + { + threw = true; + } + + BOOST_TEST_EQ(r.second, threw); + + if (!threw && !r.second) + { + const auto got_bits {std::bit_cast>(static_cast(r.first))}; + const auto expected_bits {std::bit_cast>(static_cast(expected))}; + BOOST_TEST_EQ(got_bits, expected_bits); + } + } + } +} + +template +void test_type() +{ + test_parity([](const SafeT a, const SafeT b) { return overflowing_add(a, b); }, + [](const SafeT a, const SafeT b) { return a + b; }); + test_parity([](const SafeT a, const SafeT b) { return overflowing_sub(a, b); }, + [](const SafeT a, const SafeT b) { return a - b; }); + test_parity([](const SafeT a, const SafeT b) { return overflowing_mul(a, b); }, + [](const SafeT a, const SafeT b) { return a * b; }); + test_parity([](const SafeT a, const SafeT b) { return overflowing_div(a, b); }, + [](const SafeT a, const SafeT b) { return a / b; }); +} + +// The family is constexpr and never throws, so a non-finite result is usable in +// a constant expression where the throwing operator would fail to compile (it +// would throw). The non-finite result must come from a non-finite input: +// producing inf by overflowing finite operands, or dividing by zero, is +// undefined during constant evaluation for raw floats and so cannot be +// exercised in constexpr here either. +static_assert(!overflowing_add(f32{1.0F}, f32{2.0F}).second); +static_assert(static_cast(overflowing_add(f32{1.0F}, f32{2.0F}).first) == 3.0F); +static_assert(overflowing_mul(f32{std::numeric_limits::infinity()}, f32{2.0F}).second); +static_assert(overflowing_div(f64{std::numeric_limits::infinity()}, f64{0.5}).second); +static_assert(!overflowing_sub(f64{100.0}, f64{1.0}).second); + +int main() +{ + test_type(); + test_type(); + + return boost::report_errors(); +} + +#if defined(__clang__) +# pragma clang diagnostic pop +#elif defined(__GNUC__) +# pragma GCC diagnostic pop +#endif