diff --git a/doc/modules/ROOT/pages/benchmarks.adoc b/doc/modules/ROOT/pages/benchmarks.adoc
index 0ecc79b..61452ab 100644
--- a/doc/modules/ROOT/pages/benchmarks.adoc
+++ b/doc/modules/ROOT/pages/benchmarks.adoc
@@ -56,6 +56,193 @@ The floating-point benchmarks draw values uniformly from `[1, 100]` and measure
 
 NOTE: The numbers come from shared CI runners and reflect a single run each, so treat small differences as noise. The repetition count and median reporting keep the central figure stable, but they cannot remove the variance inherent to a shared virtual machine.
 
+[#device_benchmarks]
+== Device Benchmarks (CUDA and SYCL)
+
+All of the safe types run inside CUDA and SYCL kernels (see xref:cuda.adoc[] and xref:sycl.adoc[]), and two dedicated benchmarks measure what the safety checks cost on the device itself:
+
+* `test/benchmarks/benchmark_cuda_operations.cu`, timed with CUDA events
+* `test/benchmarks/benchmark_sycl_operations.cpp`, timed with SYCL event profiling
+
+Each compares every safe type against its own basis type running the identical kernel (`boost::sn::u32` against `std::uint32_t`, `boost::sn::u128` against the emulated 128-bit basis type, and so on), so the ratio isolates the cost of the safety checks.
+As in the host benchmarks, the inputs are chosen so that no check ever fires: the timed regions never enter the error reporting path, and a `device_error_context` covering the whole run aborts the benchmark if a check does fire.
+
+[#device_benchmarks_methodology]
+=== Device Methodology
+
+Two kernel shapes are measured for every type:
+
+Streaming::
+`out[i] = a[i] op b[i]`, one checked operation per element over 4,194,304 elements.
+This is the shape real element-wise kernels have.
+On a discrete GPU it is bound by memory bandwidth, so it shows the overhead a typical kernel actually experiences.
+Integer operands are drawn from `[1, width_in_bits - 1]` with `a >= b`, and floating point operands from `[1, 100]`, mirroring the host benchmarks.
+
+Compute-bound chain::
+65,536 threads each evolve one value through a serial dependency chain held in registers (256 cycles of eight to nine checked operations), so the arithmetic cannot hide behind memory traffic and the worst case per-operation cost becomes visible.
+
+The integer chain is a fixed mixed sequence (two modulo, two addition, two division, one subtraction, and one multiplication per cycle) rather than a single repeated operation.
+That is deliberate: wrapping integer arithmetic is associative, so the optimizer can legally collapse a run of additions or multiplications by loop-invariant operands for the builtin baseline, and the comparison would then measure an optimizer artifact instead of the checks.
+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 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.
+
+[#run_device_benchmarks]
+=== How to Run the Device Benchmarks
+
+CUDA (requires the CUDA toolkit):
+
+[source, shell]
+----
+cmake -S . -B __build__ -DCMAKE_BUILD_TYPE=Release -DBUILD_TESTING=OFF \
+  -DBOOST_INCLUDE_LIBRARIES="safe_numbers" \
+  -DBOOST_SAFE_NUMBERS_BUILD_BENCHMARKS=ON \
+  -DBOOST_SAFE_NUMBERS_ENABLE_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=native
+cmake --build __build__ --target benchmark_cuda_operations
+./__build__/stage/bin/benchmark_cuda_operations
+----
+
+SYCL (requires oneAPI):
+
+[source, shell]
+----
+cmake -S . -B __build__ -DCMAKE_BUILD_TYPE=Release -DBUILD_TESTING=OFF \
+  -DCMAKE_CXX_COMPILER=icpx \
+  -DBOOST_INCLUDE_LIBRARIES="safe_numbers" \
+  -DBOOST_SAFE_NUMBERS_BUILD_BENCHMARKS=ON \
+  -DBOOST_SAFE_NUMBERS_ENABLE_SYCL=ON
+cmake --build __build__ --target benchmark_sycl_operations
+./__build__/stage/bin/benchmark_sycl_operations
+----
+
+NOTE: `icpx` compiles with `-ffp-model=fast` by default, which grants the optimizer floating point associativity; the builtin baseline chains then fold algebraically and the float comparison becomes meaningless. The build therefore pins `-ffp-model=precise` for the SYCL benchmark, so both contenders run IEEE arithmetic exactly as they do under `nvcc`. The build also splits the device code per kernel (`-fsycl-device-code-split=per_kernel`) because at `-O3` the monolithic device image of a few hundred kernels fails to JIT on the OpenCL CPU backend.
+
+[#cuda_rtx3060_benchmarks]
+=== CUDA Results: NVIDIA RTX 3060
+
+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
+|===
+
+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 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"]
+|===
+| 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
+|===
+
+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.
+Real kernels sit between the two regimes, and usually much closer to the streaming one.
+
+[#sycl_cpu_benchmarks]
+=== SYCL Results: Intel OpenCL CPU Device
+
+Run on the OpenCL CPU device of an Intel Core i9-11900K (8 cores) with oneAPI icpx 2026.0 at `-O3`, `-ffp-model=precise`, and per kernel device code split, using the default sizes.
+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
+|===
+
+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
+|===
+
+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.
+
+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.
+
 // The per-platform sections below are generated from the JSON artifacts in
 // doc/modules/ROOT/data by test/benchmarks/render_results.py. Do not edit them
 // by hand; re-run that script to refresh them.
diff --git a/doc/modules/ROOT/pages/cuda.adoc b/doc/modules/ROOT/pages/cuda.adoc
index d8a279b..1221f37 100644
--- a/doc/modules/ROOT/pages/cuda.adoc
+++ b/doc/modules/ROOT/pages/cuda.adoc
@@ -344,3 +344,7 @@ Be aware that other threads in the kernel may continue executing with incorrect
 * xref:examples.adoc#examples_cuda[CUDA Device Support] — demonstrates that all safe_numbers types and free functions work on a CUDA device.
 * 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.
+
+== 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.
diff --git a/doc/modules/ROOT/pages/sycl.adoc b/doc/modules/ROOT/pages/sycl.adoc
index 4d9b6d5..ff3624a 100644
--- a/doc/modules/ROOT/pages/sycl.adoc
+++ b/doc/modules/ROOT/pages/sycl.adoc
@@ -236,3 +236,7 @@ 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.
+
+== 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.
diff --git a/test/benchmarks/CMakeLists.txt b/test/benchmarks/CMakeLists.txt
index 9a5c99b..0cbb153 100644
--- a/test/benchmarks/CMakeLists.txt
+++ b/test/benchmarks/CMakeLists.txt
@@ -3,8 +3,56 @@
 # https://www.boost.org/LICENSE_1_0.txt
 
 # Built only when BOOST_SAFE_NUMBERS_BUILD_BENCHMARKS is ON (see the root
-# CMakeLists.txt). Google Benchmark is pulled in with FetchContent so no system
-# install is required, which keeps local and CI builds identical.
+# CMakeLists.txt). The host benchmarks use Google Benchmark, pulled in with
+# FetchContent so no system install is required. The CUDA and SYCL device
+# benchmarks time kernels with the backend's own event machinery instead, so
+# under BOOST_SAFE_NUMBERS_ENABLE_CUDA / BOOST_SAFE_NUMBERS_ENABLE_SYCL only
+# the matching device benchmark is built (mirroring test/CMakeLists.txt).
+
+if(BOOST_SAFE_NUMBERS_ENABLE_CUDA)
+
+    message(STATUS "Building the Boost.safe_numbers CUDA benchmark")
+
+    enable_language(CUDA)
+    find_package(CUDAToolkit REQUIRED)
+    set(CMAKE_CUDA_EXTENSIONS OFF)
+    set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --expt-relaxed-constexpr")
+
+    add_executable(benchmark_cuda_operations benchmark_cuda_operations.cu)
+    target_link_libraries(benchmark_cuda_operations PRIVATE Boost::safe_numbers CUDA::cudart)
+    target_compile_definitions(benchmark_cuda_operations PRIVATE BOOST_SAFE_NUMBERS_ENABLE_CUDA=1)
+    target_compile_features(benchmark_cuda_operations PRIVATE cuda_std_20)
+
+    return()
+
+elseif(BOOST_SAFE_NUMBERS_ENABLE_SYCL)
+
+    message(STATUS "Building the Boost.safe_numbers SYCL benchmark")
+
+    # Configure with -DCMAKE_CXX_COMPILER=icpx. -fsycl must be on the link line
+    # too (it bundles the device image); without it the executable segfaults at
+    # kernel submission. See test/CMakeLists.txt for the range rounding define.
+    #
+    # icpx defaults to -ffp-model=fast, which grants the optimizer floating
+    # point associativity: the builtin baseline chains then fold algebraically
+    # and the comparison is meaningless (and no longer bit-identical). Pin the
+    # precise model so both contenders run IEEE arithmetic, as nvcc does.
+    #
+    # The benchmark packs a couple of hundred kernels into one translation
+    # unit. At -O3 the resulting monolithic device image trips a JIT symbol
+    # materialization failure (old_llvm.umul.with.overflow.*) on the OpenCL
+    # CPU backend, so split the device code per kernel; the split also has to
+    # be on the link line, where the post-link tool actually runs.
+    add_executable(benchmark_sycl_operations benchmark_sycl_operations.cpp)
+    target_link_libraries(benchmark_sycl_operations PRIVATE Boost::safe_numbers sycl)
+    target_compile_definitions(benchmark_sycl_operations PRIVATE BOOST_SAFE_NUMBERS_ENABLE_SYCL=1 __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__=1)
+    target_compile_options(benchmark_sycl_operations PRIVATE -fsycl -ffp-model=precise -fsycl-device-code-split=per_kernel)
+    target_link_options(benchmark_sycl_operations PRIVATE -fsycl -fsycl-device-code-split=per_kernel)
+    target_compile_features(benchmark_sycl_operations PRIVATE cxx_std_20)
+
+    return()
+
+endif()
 
 include(FetchContent)
 
diff --git a/test/benchmarks/benchmark_cuda_operations.cu b/test/benchmarks/benchmark_cuda_operations.cu
new file mode 100644
index 0000000..20dfae6
--- /dev/null
+++ b/test/benchmarks/benchmark_cuda_operations.cu
@@ -0,0 +1,431 @@
+// Copyright 2026 Matt Borland
+// Distributed under the Boost Software License, Version 1.0.
+// https://www.boost.org/LICENSE_1_0.txt
+//
+// Measures the device-side overhead of Boost.safe_numbers on CUDA by running
+// identical kernels over the safe types and over their raw basis types.
+//
+// Two regimes are measured for every type:
+//   * streaming: out[i] = a[i] op b[i], one checked op per element. This is
+//     the memory-bound shape that real element-wise kernels have, so it shows
+//     the overhead a typical kernel actually experiences.
+//   * chain: a serial dependency chain of checked ops held in registers, the
+//     compute-bound worst case where nothing can hide behind memory traffic.
+//
+// Timing uses CUDA events around a batch of kernel launches. Every timed
+// region is verified afterwards: the safe results must match the builtin
+// results element for element, and a slice is recomputed on the host. A
+// device_error_context guards the whole run, so if any check ever fired the
+// benchmark aborts instead of silently timing the error path.
+//
+// Usage: benchmark_cuda_operations [stream_elements] [chain_cycles] [launches] [chain_threads]
+
+#ifndef BOOST_SAFE_NUMBERS_ENABLE_CUDA
+#  error "Compile with -DBOOST_SAFE_NUMBERS_ENABLE_CUDA=1 to build the CUDA benchmark"
+#endif
+
+#include <boost/safe_numbers/device_error_reporting.hpp>
+#include "device_bench.hpp"
+
+#include <cuda_runtime.h>
+
+#include <algorithm>
+#include <cstdio>
+#include <cstdlib>
+#include <exception>
+#include <vector>
+
+namespace {
+
+constexpr int block_size {256};
+
+struct bench_config
+{
+    int stream_n {1 << 22};
+    int chain_cycles {256};
+    int launches {10};
+    int chain_threads {1 << 16};
+};
+
+void cuda_check(const cudaError_t err, const char* what)
+{
+    if (err != cudaSuccess)
+    {
+        std::fprintf(stderr, "%s failed: %s\n", what, cudaGetErrorString(err));
+        std::exit(EXIT_FAILURE);
+    }
+}
+
+template <typename T>
+class device_array
+{
+public:
+    explicit device_array(const std::size_t n) : n_ {n}
+    {
+        void* ptr {};
+        cuda_check(cudaMalloc(&ptr, n * sizeof(T)), "cudaMalloc");
+        data_ = static_cast<T*>(ptr);
+    }
+
+    ~device_array()
+    {
+        cudaFree(data_);
+    }
+
+    device_array(const device_array&) = delete;
+    device_array& operator=(const device_array&) = delete;
+
+    void upload(const std::vector<T>& src)
+    {
+        cuda_check(cudaMemcpy(data_, src.data(), src.size() * sizeof(T), cudaMemcpyHostToDevice), "cudaMemcpy to device");
+    }
+
+    std::vector<T> download() const
+    {
+        std::vector<T> host(n_, device_bench::make_val<T>(0U));
+        cuda_check(cudaMemcpy(host.data(), data_, n_ * sizeof(T), cudaMemcpyDeviceToHost), "cudaMemcpy to host");
+        return host;
+    }
+
+    T* get() const
+    {
+        return data_;
+    }
+
+private:
+    T* data_ {};
+    std::size_t n_ {};
+};
+
+template <typename Op, typename T>
+__global__ void stream_kernel(const T* a, const T* b, T* out, const int n)
+{
+    const auto i {static_cast<int>(blockDim.x * blockIdx.x + threadIdx.x)};
+
+    if (i < n)
+    {
+        out[i] = Op::apply(a[i], b[i]);
+    }
+}
+
+template <typename Chain, typename T>
+__global__ void chain_kernel(const T* x0, const T* ys, const T* aux, T* out, const int cycles, const int n_threads)
+{
+    const auto t {static_cast<int>(blockDim.x * blockIdx.x + threadIdx.x)};
+
+    if (t < n_threads)
+    {
+        const auto base {static_cast<long long>(t) * device_bench::chain_operands};
+        const T y[device_bench::chain_operands] {ys[base], ys[base + 1], ys[base + 2], ys[base + 3],
+                                                 ys[base + 4], ys[base + 5], ys[base + 6], ys[base + 7]};
+        const T a {aux[t]};
+        T x {x0[t]};
+
+        for (int r {0}; r < cycles; ++r)
+        {
+            x = Chain::cycle(x, y, a);
+        }
+
+        out[t] = x;
+    }
+}
+
+// Returns the average time of one launch in nanoseconds
+template <typename F>
+double time_launches_ns(F launch, const int launches)
+{
+    launch();
+    launch();
+    cuda_check(cudaGetLastError(), "warmup launch");
+    cuda_check(cudaDeviceSynchronize(), "warmup synchronize");
+
+    cudaEvent_t start {};
+    cudaEvent_t stop {};
+    cuda_check(cudaEventCreate(&start), "cudaEventCreate");
+    cuda_check(cudaEventCreate(&stop), "cudaEventCreate");
+
+    cuda_check(cudaEventRecord(start), "cudaEventRecord");
+    for (int i {0}; i < launches; ++i)
+    {
+        launch();
+    }
+    cuda_check(cudaEventRecord(stop), "cudaEventRecord");
+    cuda_check(cudaGetLastError(), "timed launch");
+    cuda_check(cudaEventSynchronize(stop), "cudaEventSynchronize");
+
+    auto ms {0.0F};
+    cuda_check(cudaEventElapsedTime(&ms, start, stop), "cudaEventElapsedTime");
+
+    cudaEventDestroy(start);
+    cudaEventDestroy(stop);
+
+    return static_cast<double>(ms) * 1.0e6 / launches;
+}
+
+template <typename Op, typename TB, typename TS>
+int run_stream_op(boost::safe_numbers::device_error_context& ctx,
+                  const device_bench::stream_inputs<TB, TS>& host_in,
+                  const device_array<TB>& a_b, const device_array<TB>& b_b, device_array<TB>& out_b,
+                  const device_array<TS>& a_s, const device_array<TS>& b_s, device_array<TS>& out_s,
+                  const bench_config& cfg)
+{
+    const auto n {cfg.stream_n};
+    const auto blocks {(n + block_size - 1) / block_size};
+
+    const auto builtin_ns {time_launches_ns([&]()
+    {
+        stream_kernel<Op><<<blocks, block_size>>>(a_b.get(), b_b.get(), out_b.get(), n);
+    }, cfg.launches)};
+    ctx.synchronize();
+
+    const auto safe_ns {time_launches_ns([&]()
+    {
+        stream_kernel<Op><<<blocks, block_size>>>(a_s.get(), b_s.get(), out_s.get(), n);
+    }, cfg.launches)};
+    ctx.synchronize();
+
+    const auto host_b {out_b.download()};
+    const auto host_s {out_s.download()};
+
+    auto bad {device_bench::count_mismatches(host_b, host_s)};
+
+    const auto spot {std::min(n, 1024)};
+    for (int i {0}; i < spot; ++i)
+    {
+        if (host_b[i] != Op::apply(host_in.a.builtin_vals[static_cast<std::size_t>(i)],
+                                   host_in.b.builtin_vals[static_cast<std::size_t>(i)]))
+        {
+            ++bad;
+        }
+    }
+
+    return device_bench::print_row(Op::name, builtin_ns / n, safe_ns / n, bad);
+}
+
+template <template <typename> class Chain, typename TB, typename TS>
+int run_chain(boost::safe_numbers::device_error_context& ctx,
+              const device_bench::chain_inputs<TB, TS>& host_in,
+              const bench_config& cfg)
+{
+    const auto threads {cfg.chain_threads};
+    const auto blocks {(threads + block_size - 1) / block_size};
+
+    device_array<TB> x0_b {static_cast<std::size_t>(threads)};
+    device_array<TB> ys_b {static_cast<std::size_t>(threads) * device_bench::chain_operands};
+    device_array<TB> aux_b {static_cast<std::size_t>(threads)};
+    device_array<TB> out_b {static_cast<std::size_t>(threads)};
+    x0_b.upload(host_in.x0.builtin_vals);
+    ys_b.upload(host_in.ys.builtin_vals);
+    aux_b.upload(host_in.aux.builtin_vals);
+
+    device_array<TS> x0_s {static_cast<std::size_t>(threads)};
+    device_array<TS> ys_s {static_cast<std::size_t>(threads) * device_bench::chain_operands};
+    device_array<TS> aux_s {static_cast<std::size_t>(threads)};
+    device_array<TS> out_s {static_cast<std::size_t>(threads)};
+    x0_s.upload(host_in.x0.safe_vals);
+    ys_s.upload(host_in.ys.safe_vals);
+    aux_s.upload(host_in.aux.safe_vals);
+
+    const auto builtin_ns {time_launches_ns([&]()
+    {
+        chain_kernel<Chain<TB>><<<blocks, block_size>>>(x0_b.get(), ys_b.get(), aux_b.get(), out_b.get(),
+                                                        cfg.chain_cycles, threads);
+    }, cfg.launches)};
+    ctx.synchronize();
+
+    const auto safe_ns {time_launches_ns([&]()
+    {
+        chain_kernel<Chain<TS>><<<blocks, block_size>>>(x0_s.get(), ys_s.get(), aux_s.get(), out_s.get(),
+                                                        cfg.chain_cycles, threads);
+    }, cfg.launches)};
+    ctx.synchronize();
+
+    const auto host_b {out_b.download()};
+    const auto host_s {out_s.download()};
+
+    auto bad {device_bench::count_mismatches(host_b, host_s)};
+
+    // Recompute a slice of the chains on the host to catch shared miscompiles
+    const auto spot_threads {std::min(threads, 128)};
+    for (int t {0}; t < spot_threads; ++t)
+    {
+        const auto base {static_cast<std::size_t>(t) * device_bench::chain_operands};
+        TB y[device_bench::chain_operands] {};
+        for (int k {0}; k < device_bench::chain_operands; ++k)
+        {
+            y[k] = host_in.ys.builtin_vals[base + static_cast<std::size_t>(k)];
+        }
+
+        auto x {host_in.x0.builtin_vals[static_cast<std::size_t>(t)]};
+        const auto a {host_in.aux.builtin_vals[static_cast<std::size_t>(t)]};
+        for (int r {0}; r < cfg.chain_cycles; ++r)
+        {
+            x = Chain<TB>::cycle(x, y, a);
+        }
+
+        if (host_b[static_cast<std::size_t>(t)] != x)
+        {
+            ++bad;
+        }
+    }
+
+    const auto ops {static_cast<double>(threads) * cfg.chain_cycles * Chain<TB>::ops_per_cycle};
+    return device_bench::print_row(Chain<TB>::name, builtin_ns / ops, safe_ns / ops, bad);
+}
+
+template <typename TS>
+int run_int_family(const char* name, boost::safe_numbers::device_error_context& ctx, const bench_config& cfg)
+{
+    using TB = typename TS::basis_type;
+
+    auto failures {0};
+
+    {
+        const auto in {device_bench::gen_stream_int<TB, TS>(cfg.stream_n)};
+
+        device_array<TB> a_b {static_cast<std::size_t>(cfg.stream_n)};
+        device_array<TB> b_b {static_cast<std::size_t>(cfg.stream_n)};
+        device_array<TB> out_b {static_cast<std::size_t>(cfg.stream_n)};
+        a_b.upload(in.a.builtin_vals);
+        b_b.upload(in.b.builtin_vals);
+
+        device_array<TS> a_s {static_cast<std::size_t>(cfg.stream_n)};
+        device_array<TS> b_s {static_cast<std::size_t>(cfg.stream_n)};
+        device_array<TS> out_s {static_cast<std::size_t>(cfg.stream_n)};
+        a_s.upload(in.a.safe_vals);
+        b_s.upload(in.b.safe_vals);
+
+        device_bench::print_table_header(name, "streaming, one checked op per element (memory-bound)");
+        failures += run_stream_op<device_bench::op_add>(ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg);
+        failures += run_stream_op<device_bench::op_sub>(ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg);
+        failures += run_stream_op<device_bench::op_mul>(ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg);
+        failures += run_stream_op<device_bench::op_div>(ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg);
+        failures += run_stream_op<device_bench::op_mod>(ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg);
+    }
+
+    {
+        const auto in {device_bench::gen_int_chain<TB, TS>(cfg.chain_threads)};
+
+        device_bench::print_table_header(name, "register chain of checked ops (compute-bound)");
+        failures += run_chain<device_bench::int_mixed_chain>(ctx, in, cfg);
+    }
+
+    return failures;
+}
+
+template <typename TS>
+int run_real_family(const char* name, boost::safe_numbers::device_error_context& ctx, const bench_config& cfg)
+{
+    using TB = typename TS::basis_type;
+
+    auto failures {0};
+
+    {
+        const auto in {device_bench::gen_stream_real<TB, TS>(cfg.stream_n)};
+
+        device_array<TB> a_b {static_cast<std::size_t>(cfg.stream_n)};
+        device_array<TB> b_b {static_cast<std::size_t>(cfg.stream_n)};
+        device_array<TB> out_b {static_cast<std::size_t>(cfg.stream_n)};
+        a_b.upload(in.a.builtin_vals);
+        b_b.upload(in.b.builtin_vals);
+
+        device_array<TS> a_s {static_cast<std::size_t>(cfg.stream_n)};
+        device_array<TS> b_s {static_cast<std::size_t>(cfg.stream_n)};
+        device_array<TS> out_s {static_cast<std::size_t>(cfg.stream_n)};
+        a_s.upload(in.a.safe_vals);
+        b_s.upload(in.b.safe_vals);
+
+        device_bench::print_table_header(name, "streaming, one checked op per element (memory-bound)");
+        failures += run_stream_op<device_bench::op_add>(ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg);
+        failures += run_stream_op<device_bench::op_sub>(ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg);
+        failures += run_stream_op<device_bench::op_mul>(ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg);
+        failures += run_stream_op<device_bench::op_div>(ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg);
+    }
+
+    {
+        device_bench::print_table_header(name, "register chain of checked ops (compute-bound)");
+
+        const auto add_in {device_bench::gen_real_chain<TB, TS>(cfg.chain_threads, device_bench::real_chain_kind::sum)};
+        failures += run_chain<device_bench::real_add_chain>(ctx, add_in, cfg);
+        failures += run_chain<device_bench::real_sub_chain>(ctx, add_in, cfg);
+
+        const auto mul_in {device_bench::gen_real_chain<TB, TS>(cfg.chain_threads, device_bench::real_chain_kind::product)};
+        failures += run_chain<device_bench::real_mul_chain>(ctx, mul_in, cfg);
+
+        const auto div_in {device_bench::gen_real_chain<TB, TS>(cfg.chain_threads, device_bench::real_chain_kind::divisor)};
+        failures += run_chain<device_bench::real_div_chain>(ctx, div_in, cfg);
+    }
+
+    return failures;
+}
+
+} // namespace
+
+int main(int argc, char** argv)
+{
+    bench_config cfg {};
+    if (argc > 1)
+    {
+        cfg.stream_n = std::atoi(argv[1]);
+    }
+    if (argc > 2)
+    {
+        cfg.chain_cycles = std::atoi(argv[2]);
+    }
+    if (argc > 3)
+    {
+        cfg.launches = std::atoi(argv[3]);
+    }
+    if (argc > 4)
+    {
+        cfg.chain_threads = std::atoi(argv[4]);
+    }
+    if (cfg.stream_n <= 0 || cfg.chain_cycles <= 0 || cfg.launches <= 0 || cfg.chain_threads <= 0)
+    {
+        std::fprintf(stderr, "usage: %s [stream_elements] [chain_cycles] [launches] [chain_threads]\n", argv[0]);
+        return EXIT_FAILURE;
+    }
+
+    cudaDeviceProp prop {};
+    cuda_check(cudaGetDeviceProperties(&prop, 0), "cudaGetDeviceProperties");
+    std::printf("Boost.safe_numbers CUDA benchmark\n");
+    std::printf("Device: %s (sm_%d%d, %d SMs)\n", prop.name, prop.major, prop.minor, prop.multiProcessorCount);
+    std::printf("Streaming elements: %d, chain threads: %d, chain cycles: %d, launches per timing: %d\n",
+                cfg.stream_n, cfg.chain_threads, cfg.chain_cycles, cfg.launches);
+
+    auto failures {0};
+
+    try
+    {
+        boost::safe_numbers::device_error_context ctx;
+
+        failures += run_int_family<boost::safe_numbers::u8>("u8", ctx, cfg);
+        failures += run_int_family<boost::safe_numbers::u16>("u16", ctx, cfg);
+        failures += run_int_family<boost::safe_numbers::u32>("u32", ctx, cfg);
+        failures += run_int_family<boost::safe_numbers::u64>("u64", ctx, cfg);
+        failures += run_int_family<boost::safe_numbers::u128>("u128", ctx, cfg);
+
+        failures += run_int_family<boost::safe_numbers::i8>("i8", ctx, cfg);
+        failures += run_int_family<boost::safe_numbers::i16>("i16", ctx, cfg);
+        failures += run_int_family<boost::safe_numbers::i32>("i32", ctx, cfg);
+        failures += run_int_family<boost::safe_numbers::i64>("i64", ctx, cfg);
+        failures += run_int_family<boost::safe_numbers::i128>("i128", ctx, cfg);
+
+        failures += run_real_family<boost::safe_numbers::f32>("f32", ctx, cfg);
+        failures += run_real_family<boost::safe_numbers::f64>("f64", ctx, cfg);
+    }
+    catch (const std::exception& e)
+    {
+        std::fprintf(stderr, "\nUnexpected error (a safety check fired inside a timed region): %s\n", e.what());
+        return EXIT_FAILURE;
+    }
+
+    if (failures != 0)
+    {
+        std::printf("\n%d benchmark(s) FAILED verification\n", failures);
+        return EXIT_FAILURE;
+    }
+
+    std::printf("\nAll benchmarks verified\n");
+    return EXIT_SUCCESS;
+}
diff --git a/test/benchmarks/benchmark_sycl_operations.cpp b/test/benchmarks/benchmark_sycl_operations.cpp
new file mode 100644
index 0000000..d21a422
--- /dev/null
+++ b/test/benchmarks/benchmark_sycl_operations.cpp
@@ -0,0 +1,411 @@
+// Copyright 2026 Matt Borland
+// Distributed under the Boost Software License, Version 1.0.
+// https://www.boost.org/LICENSE_1_0.txt
+//
+// Measures the device-side overhead of Boost.safe_numbers on SYCL by running
+// identical kernels over the safe types and over their raw basis types.
+// The structure mirrors benchmark_cuda_operations.cu; see that file and
+// device_bench.hpp for the methodology.
+//
+// Timing sums per-kernel device execution time from SYCL event profiling, so
+// submission gaps are excluded. Every timed region is verified afterwards and
+// a device_error_context guards the whole run, so if any check ever fired the
+// benchmark aborts instead of silently timing the error path.
+//
+// Build with icpx: -fsycl (compile and link) -DBOOST_SAFE_NUMBERS_ENABLE_SYCL=1
+//
+// Usage: benchmark_sycl_operations [stream_elements] [chain_cycles] [launches] [chain_threads]
+
+#include <sycl/sycl.hpp>
+
+#ifndef BOOST_SAFE_NUMBERS_ENABLE_SYCL
+#  error "Compile with -DBOOST_SAFE_NUMBERS_ENABLE_SYCL=1 to build the SYCL benchmark"
+#endif
+
+#include <boost/safe_numbers/device_error_reporting.hpp>
+#include "device_bench.hpp"
+
+#include <algorithm>
+#include <cstdint>
+#include <cstdio>
+#include <cstdlib>
+#include <exception>
+#include <stdexcept>
+#include <string>
+#include <vector>
+
+namespace {
+
+struct bench_config
+{
+    int stream_n {1 << 22};
+    int chain_cycles {256};
+    int launches {10};
+    int chain_threads {1 << 16};
+};
+
+template <typename T>
+class usm_array
+{
+public:
+    usm_array(sycl::queue& q, const std::size_t n) : q_ {&q}, n_ {n}, data_ {sycl::malloc_device<T>(n, q)}
+    {
+        if (data_ == nullptr)
+        {
+            throw std::runtime_error("sycl::malloc_device failed");
+        }
+    }
+
+    ~usm_array()
+    {
+        sycl::free(data_, *q_);
+    }
+
+    usm_array(const usm_array&) = delete;
+    usm_array& operator=(const usm_array&) = delete;
+
+    void upload(const std::vector<T>& src)
+    {
+        q_->memcpy(data_, src.data(), src.size() * sizeof(T)).wait();
+    }
+
+    std::vector<T> download() const
+    {
+        std::vector<T> host(n_, device_bench::make_val<T>(0U));
+        q_->memcpy(host.data(), data_, n_ * sizeof(T)).wait();
+        return host;
+    }
+
+    T* get() const
+    {
+        return data_;
+    }
+
+private:
+    sycl::queue* q_ {};
+    std::size_t n_ {};
+    T* data_ {};
+};
+
+// Returns the average device execution time of one launch in nanoseconds,
+// summed from event profiling so host-side submission gaps are excluded
+template <typename F>
+double time_launches_ns(sycl::queue& q, F launch, const int launches)
+{
+    launch().wait();
+    launch().wait();
+
+    std::vector<sycl::event> events;
+    events.reserve(static_cast<std::size_t>(launches));
+    for (int i {0}; i < launches; ++i)
+    {
+        events.push_back(launch());
+    }
+    q.wait();
+
+    auto total_ns {std::uint64_t {0}};
+    for (auto& ev : events)
+    {
+        const auto begin {ev.get_profiling_info<sycl::info::event_profiling::command_start>()};
+        const auto end {ev.get_profiling_info<sycl::info::event_profiling::command_end>()};
+        total_ns += end - begin;
+    }
+
+    return static_cast<double>(total_ns) / launches;
+}
+
+template <typename Op, typename T>
+sycl::event launch_stream(sycl::queue& q, const T* a, const T* b, T* out, const int n)
+{
+    return q.parallel_for(sycl::range<1>(static_cast<std::size_t>(n)), [=](sycl::id<1> idx)
+    {
+        const auto i {idx[0]};
+        out[i] = Op::apply(a[i], b[i]);
+    });
+}
+
+template <typename Chain, typename T>
+sycl::event launch_chain(sycl::queue& q, const T* x0, const T* ys, const T* aux, T* out, const int cycles, const int threads)
+{
+    return q.parallel_for(sycl::range<1>(static_cast<std::size_t>(threads)), [=](sycl::id<1> idx)
+    {
+        const auto t {idx[0]};
+        const auto base {t * device_bench::chain_operands};
+        const T y[device_bench::chain_operands] {ys[base], ys[base + 1], ys[base + 2], ys[base + 3],
+                                                 ys[base + 4], ys[base + 5], ys[base + 6], ys[base + 7]};
+        const T a {aux[t]};
+        T x {x0[t]};
+
+        for (int r {0}; r < cycles; ++r)
+        {
+            x = Chain::cycle(x, y, a);
+        }
+
+        out[t] = x;
+    });
+}
+
+template <typename Op, typename TB, typename TS>
+int run_stream_op(sycl::queue& q, boost::safe_numbers::device_error_context& ctx,
+                  const device_bench::stream_inputs<TB, TS>& host_in,
+                  const usm_array<TB>& a_b, const usm_array<TB>& b_b, usm_array<TB>& out_b,
+                  const usm_array<TS>& a_s, const usm_array<TS>& b_s, usm_array<TS>& out_s,
+                  const bench_config& cfg)
+{
+    const auto n {cfg.stream_n};
+
+    const auto builtin_ns {time_launches_ns(q, [&]()
+    {
+        return launch_stream<Op>(q, a_b.get(), b_b.get(), out_b.get(), n);
+    }, cfg.launches)};
+    ctx.synchronize();
+
+    const auto safe_ns {time_launches_ns(q, [&]()
+    {
+        return launch_stream<Op>(q, a_s.get(), b_s.get(), out_s.get(), n);
+    }, cfg.launches)};
+    ctx.synchronize();
+
+    const auto host_b {out_b.download()};
+    const auto host_s {out_s.download()};
+
+    auto bad {device_bench::count_mismatches(host_b, host_s)};
+
+    const auto spot {std::min(n, 1024)};
+    for (int i {0}; i < spot; ++i)
+    {
+        if (host_b[static_cast<std::size_t>(i)] != Op::apply(host_in.a.builtin_vals[static_cast<std::size_t>(i)],
+                                                             host_in.b.builtin_vals[static_cast<std::size_t>(i)]))
+        {
+            ++bad;
+        }
+    }
+
+    return device_bench::print_row(Op::name, builtin_ns / n, safe_ns / n, bad);
+}
+
+template <template <typename> class Chain, typename TB, typename TS>
+int run_chain(sycl::queue& q, boost::safe_numbers::device_error_context& ctx,
+              const device_bench::chain_inputs<TB, TS>& host_in,
+              const bench_config& cfg)
+{
+    const auto threads {cfg.chain_threads};
+
+    usm_array<TB> x0_b {q, static_cast<std::size_t>(threads)};
+    usm_array<TB> ys_b {q, static_cast<std::size_t>(threads) * device_bench::chain_operands};
+    usm_array<TB> aux_b {q, static_cast<std::size_t>(threads)};
+    usm_array<TB> out_b {q, static_cast<std::size_t>(threads)};
+    x0_b.upload(host_in.x0.builtin_vals);
+    ys_b.upload(host_in.ys.builtin_vals);
+    aux_b.upload(host_in.aux.builtin_vals);
+
+    usm_array<TS> x0_s {q, static_cast<std::size_t>(threads)};
+    usm_array<TS> ys_s {q, static_cast<std::size_t>(threads) * device_bench::chain_operands};
+    usm_array<TS> aux_s {q, static_cast<std::size_t>(threads)};
+    usm_array<TS> out_s {q, static_cast<std::size_t>(threads)};
+    x0_s.upload(host_in.x0.safe_vals);
+    ys_s.upload(host_in.ys.safe_vals);
+    aux_s.upload(host_in.aux.safe_vals);
+
+    const auto builtin_ns {time_launches_ns(q, [&]()
+    {
+        return launch_chain<Chain<TB>>(q, x0_b.get(), ys_b.get(), aux_b.get(), out_b.get(), cfg.chain_cycles, threads);
+    }, cfg.launches)};
+    ctx.synchronize();
+
+    const auto safe_ns {time_launches_ns(q, [&]()
+    {
+        return launch_chain<Chain<TS>>(q, x0_s.get(), ys_s.get(), aux_s.get(), out_s.get(), cfg.chain_cycles, threads);
+    }, cfg.launches)};
+    ctx.synchronize();
+
+    const auto host_b {out_b.download()};
+    const auto host_s {out_s.download()};
+
+    auto bad {device_bench::count_mismatches(host_b, host_s)};
+
+    // Recompute a slice of the chains on the host to catch shared miscompiles
+    const auto spot_threads {std::min(threads, 128)};
+    for (int t {0}; t < spot_threads; ++t)
+    {
+        const auto base {static_cast<std::size_t>(t) * device_bench::chain_operands};
+        TB y[device_bench::chain_operands] {};
+        for (int k {0}; k < device_bench::chain_operands; ++k)
+        {
+            y[k] = host_in.ys.builtin_vals[base + static_cast<std::size_t>(k)];
+        }
+
+        auto x {host_in.x0.builtin_vals[static_cast<std::size_t>(t)]};
+        const auto a {host_in.aux.builtin_vals[static_cast<std::size_t>(t)]};
+        for (int r {0}; r < cfg.chain_cycles; ++r)
+        {
+            x = Chain<TB>::cycle(x, y, a);
+        }
+
+        if (host_b[static_cast<std::size_t>(t)] != x)
+        {
+            ++bad;
+        }
+    }
+
+    const auto ops {static_cast<double>(threads) * cfg.chain_cycles * Chain<TB>::ops_per_cycle};
+    return device_bench::print_row(Chain<TB>::name, builtin_ns / ops, safe_ns / ops, bad);
+}
+
+template <typename TS>
+int run_int_family(const char* name, sycl::queue& q, boost::safe_numbers::device_error_context& ctx, const bench_config& cfg)
+{
+    using TB = typename TS::basis_type;
+
+    auto failures {0};
+
+    {
+        const auto in {device_bench::gen_stream_int<TB, TS>(cfg.stream_n)};
+
+        usm_array<TB> a_b {q, static_cast<std::size_t>(cfg.stream_n)};
+        usm_array<TB> b_b {q, static_cast<std::size_t>(cfg.stream_n)};
+        usm_array<TB> out_b {q, static_cast<std::size_t>(cfg.stream_n)};
+        a_b.upload(in.a.builtin_vals);
+        b_b.upload(in.b.builtin_vals);
+
+        usm_array<TS> a_s {q, static_cast<std::size_t>(cfg.stream_n)};
+        usm_array<TS> b_s {q, static_cast<std::size_t>(cfg.stream_n)};
+        usm_array<TS> out_s {q, static_cast<std::size_t>(cfg.stream_n)};
+        a_s.upload(in.a.safe_vals);
+        b_s.upload(in.b.safe_vals);
+
+        device_bench::print_table_header(name, "streaming, one checked op per element (memory-bound)");
+        failures += run_stream_op<device_bench::op_add>(q, ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg);
+        failures += run_stream_op<device_bench::op_sub>(q, ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg);
+        failures += run_stream_op<device_bench::op_mul>(q, ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg);
+        failures += run_stream_op<device_bench::op_div>(q, ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg);
+        failures += run_stream_op<device_bench::op_mod>(q, ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg);
+    }
+
+    {
+        const auto in {device_bench::gen_int_chain<TB, TS>(cfg.chain_threads)};
+
+        device_bench::print_table_header(name, "register chain of checked ops (compute-bound)");
+        failures += run_chain<device_bench::int_mixed_chain>(q, ctx, in, cfg);
+    }
+
+    return failures;
+}
+
+template <typename TS>
+int run_real_family(const char* name, sycl::queue& q, boost::safe_numbers::device_error_context& ctx, const bench_config& cfg)
+{
+    using TB = typename TS::basis_type;
+
+    auto failures {0};
+
+    {
+        const auto in {device_bench::gen_stream_real<TB, TS>(cfg.stream_n)};
+
+        usm_array<TB> a_b {q, static_cast<std::size_t>(cfg.stream_n)};
+        usm_array<TB> b_b {q, static_cast<std::size_t>(cfg.stream_n)};
+        usm_array<TB> out_b {q, static_cast<std::size_t>(cfg.stream_n)};
+        a_b.upload(in.a.builtin_vals);
+        b_b.upload(in.b.builtin_vals);
+
+        usm_array<TS> a_s {q, static_cast<std::size_t>(cfg.stream_n)};
+        usm_array<TS> b_s {q, static_cast<std::size_t>(cfg.stream_n)};
+        usm_array<TS> out_s {q, static_cast<std::size_t>(cfg.stream_n)};
+        a_s.upload(in.a.safe_vals);
+        b_s.upload(in.b.safe_vals);
+
+        device_bench::print_table_header(name, "streaming, one checked op per element (memory-bound)");
+        failures += run_stream_op<device_bench::op_add>(q, ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg);
+        failures += run_stream_op<device_bench::op_sub>(q, ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg);
+        failures += run_stream_op<device_bench::op_mul>(q, ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg);
+        failures += run_stream_op<device_bench::op_div>(q, ctx, in, a_b, b_b, out_b, a_s, b_s, out_s, cfg);
+    }
+
+    {
+        device_bench::print_table_header(name, "register chain of checked ops (compute-bound)");
+
+        const auto add_in {device_bench::gen_real_chain<TB, TS>(cfg.chain_threads, device_bench::real_chain_kind::sum)};
+        failures += run_chain<device_bench::real_add_chain>(q, ctx, add_in, cfg);
+        failures += run_chain<device_bench::real_sub_chain>(q, ctx, add_in, cfg);
+
+        const auto mul_in {device_bench::gen_real_chain<TB, TS>(cfg.chain_threads, device_bench::real_chain_kind::product)};
+        failures += run_chain<device_bench::real_mul_chain>(q, ctx, mul_in, cfg);
+
+        const auto div_in {device_bench::gen_real_chain<TB, TS>(cfg.chain_threads, device_bench::real_chain_kind::divisor)};
+        failures += run_chain<device_bench::real_div_chain>(q, ctx, div_in, cfg);
+    }
+
+    return failures;
+}
+
+} // namespace
+
+int main(int argc, char** argv)
+{
+    bench_config cfg {};
+    if (argc > 1)
+    {
+        cfg.stream_n = std::atoi(argv[1]);
+    }
+    if (argc > 2)
+    {
+        cfg.chain_cycles = std::atoi(argv[2]);
+    }
+    if (argc > 3)
+    {
+        cfg.launches = std::atoi(argv[3]);
+    }
+    if (argc > 4)
+    {
+        cfg.chain_threads = std::atoi(argv[4]);
+    }
+    if (cfg.stream_n <= 0 || cfg.chain_cycles <= 0 || cfg.launches <= 0 || cfg.chain_threads <= 0)
+    {
+        std::fprintf(stderr, "usage: %s [stream_elements] [chain_cycles] [launches] [chain_threads]\n", argv[0]);
+        return EXIT_FAILURE;
+    }
+
+    auto failures {0};
+
+    try
+    {
+        sycl::queue q {sycl::property_list{sycl::property::queue::in_order{}, sycl::property::queue::enable_profiling{}}};
+
+        const auto device_name {q.get_device().get_info<sycl::info::device::name>()};
+        std::printf("Boost.safe_numbers SYCL benchmark\n");
+        std::printf("Device: %s\n", device_name.c_str());
+        std::printf("Streaming elements: %d, chain threads: %d, chain cycles: %d, launches per timing: %d\n",
+                    cfg.stream_n, cfg.chain_threads, cfg.chain_cycles, cfg.launches);
+
+        boost::safe_numbers::device_error_context ctx {q};
+
+        failures += run_int_family<boost::safe_numbers::u8>("u8", q, ctx, cfg);
+        failures += run_int_family<boost::safe_numbers::u16>("u16", q, ctx, cfg);
+        failures += run_int_family<boost::safe_numbers::u32>("u32", q, ctx, cfg);
+        failures += run_int_family<boost::safe_numbers::u64>("u64", q, ctx, cfg);
+        failures += run_int_family<boost::safe_numbers::u128>("u128", q, ctx, cfg);
+
+        failures += run_int_family<boost::safe_numbers::i8>("i8", q, ctx, cfg);
+        failures += run_int_family<boost::safe_numbers::i16>("i16", q, ctx, cfg);
+        failures += run_int_family<boost::safe_numbers::i32>("i32", q, ctx, cfg);
+        failures += run_int_family<boost::safe_numbers::i64>("i64", q, ctx, cfg);
+        failures += run_int_family<boost::safe_numbers::i128>("i128", q, ctx, cfg);
+
+        failures += run_real_family<boost::safe_numbers::f32>("f32", q, ctx, cfg);
+        failures += run_real_family<boost::safe_numbers::f64>("f64", q, ctx, cfg);
+    }
+    catch (const std::exception& e)
+    {
+        std::fprintf(stderr, "\nUnexpected error (a safety check fired inside a timed region): %s\n", e.what());
+        return EXIT_FAILURE;
+    }
+
+    if (failures != 0)
+    {
+        std::printf("\n%d benchmark(s) FAILED verification\n", failures);
+        return EXIT_FAILURE;
+    }
+
+    std::printf("\nAll benchmarks verified\n");
+    return EXIT_SUCCESS;
+}
diff --git a/test/benchmarks/device_bench.hpp b/test/benchmarks/device_bench.hpp
new file mode 100644
index 0000000..9a0f3fd
--- /dev/null
+++ b/test/benchmarks/device_bench.hpp
@@ -0,0 +1,475 @@
+// Copyright 2026 Matt Borland
+// Distributed under the Boost Software License, Version 1.0.
+// https://www.boost.org/LICENSE_1_0.txt
+//
+// Shared pieces for the CUDA and SYCL device benchmarks: the benchmarked
+// operations, the compute-bound chain recurrences, host-side input generation,
+// and result-table printing. The backend files own memory management, kernel
+// launch, and timing.
+//
+// Every benchmark compares the safe type against its own basis type running
+// the identical kernel, so the ratio isolates the cost of the safety checks.
+// Input ranges are chosen so that no check ever fires, even for the 8-bit
+// types; the error path is never part of a timed region.
+
+#ifndef BOOST_SAFE_NUMBERS_BENCHMARKS_DEVICE_BENCH_HPP
+#define BOOST_SAFE_NUMBERS_BENCHMARKS_DEVICE_BENCH_HPP
+
+#include <boost/safe_numbers/unsigned_integers.hpp>
+#include <boost/safe_numbers/signed_integers.hpp>
+#include <boost/safe_numbers/floats.hpp>
+#include <cstdint>
+#include <cstdio>
+#include <random>
+#include <utility>
+#include <vector>
+
+namespace device_bench {
+
+// Number of per-thread operands carried by every compute-bound chain
+inline constexpr int chain_operands {8};
+
+// ---------------------------------------------------------------------------
+// Value helpers that work for both builtin types and safe_numbers types
+// ---------------------------------------------------------------------------
+
+template <typename T, typename U>
+T make_val(const U v)
+{
+    if constexpr (requires { typename T::basis_type; })
+    {
+        return T{static_cast<typename T::basis_type>(v)};
+    }
+    else
+    {
+        return static_cast<T>(v);
+    }
+}
+
+// Convert a benchmarked value back to the builtin basis type for verification
+template <typename TB, typename TS>
+TB as_builtin(const TS v)
+{
+    if constexpr (requires { typename TS::basis_type; })
+    {
+        return static_cast<TB>(static_cast<typename TS::basis_type>(v));
+    }
+    else
+    {
+        return static_cast<TB>(v);
+    }
+}
+
+// ---------------------------------------------------------------------------
+// Streaming ops: out[i] = op(a[i], b[i]), one checked operation per element.
+// This is the memory-bound regime that real element-wise kernels live in.
+// ---------------------------------------------------------------------------
+
+struct op_add
+{
+    static constexpr const char* name {"add"};
+
+    template <typename T>
+    BOOST_SAFE_NUMBERS_HOST_DEVICE static T apply(const T a, const T b)
+    {
+        return static_cast<T>(a + b);
+    }
+};
+
+struct op_sub
+{
+    static constexpr const char* name {"sub"};
+
+    template <typename T>
+    BOOST_SAFE_NUMBERS_HOST_DEVICE static T apply(const T a, const T b)
+    {
+        return static_cast<T>(a - b);
+    }
+};
+
+struct op_mul
+{
+    static constexpr const char* name {"mul"};
+
+    template <typename T>
+    BOOST_SAFE_NUMBERS_HOST_DEVICE static T apply(const T a, const T b)
+    {
+        return static_cast<T>(a * b);
+    }
+};
+
+struct op_div
+{
+    static constexpr const char* name {"div"};
+
+    template <typename T>
+    BOOST_SAFE_NUMBERS_HOST_DEVICE static T apply(const T a, const T b)
+    {
+        return static_cast<T>(a / b);
+    }
+};
+
+struct op_mod
+{
+    static constexpr const char* name {"mod"};
+
+    template <typename T>
+    BOOST_SAFE_NUMBERS_HOST_DEVICE static T apply(const T a, const T b)
+    {
+        return static_cast<T>(a % b);
+    }
+};
+
+// ---------------------------------------------------------------------------
+// Compute-bound chains
+//
+// Each thread evolves one value through a serial dependency chain held in
+// registers, so the arithmetic cannot hide behind memory traffic and the
+// per-operation check cost becomes visible.
+//
+// The integer cycle deliberately interleaves division and modulo between the
+// add/sub/mul steps: a run of adds or muls with loop-invariant operands is
+// affine, and the optimizer may legally collapse it for the builtin baseline
+// (wrapping arithmetic is associative), which would corrupt the comparison.
+// Division and modulo are not reassociable, so every step here must execute.
+//
+// The floating-point cycles are per-operation: IEEE arithmetic is not
+// associative, so the compiler cannot fold repeated ops, and each cycle ends
+// with one restore step of the inverse kind that returns the value to its
+// starting range.
+// ---------------------------------------------------------------------------
+
+// Operand slots (generated by gen_int_chain):
+//   y[0] in [11, 13]  y[1] in [8, 15]  y[2] in [1, 2]  y[3] in [1, 4]
+//   y[4] in [1, 4]    y[5] in [5, 9]   y[6] in [1, 7]  y[7] in [1, 2]
+// Value bounds are annotated per step; the maximum of 104 fits every type
+// down to i8, divisors are never zero, and values are never negative.
+template <typename T>
+struct int_mixed_chain
+{
+    static constexpr int ops_per_cycle {8};
+    static constexpr const char* name {"mixed"};
+
+    BOOST_SAFE_NUMBERS_HOST_DEVICE static T cycle(T x, const T* y, const T aux)
+    {
+        static_cast<void>(aux);
+        x = static_cast<T>(x % y[0]); // [0, 12]
+        x = static_cast<T>(x + y[1]); // [8, 27]
+        x = static_cast<T>(x / y[2]); // [4, 27]
+        x = static_cast<T>(x - y[3]); // [0, 26]
+        x = static_cast<T>(x * y[4]); // [0, 104]
+        x = static_cast<T>(x % y[5]); // [0, 8]
+        x = static_cast<T>(x + y[6]); // [1, 15]
+        x = static_cast<T>(x / y[7]); // [0, 15]
+        return x;
+    }
+};
+
+// 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 <typename T>
+struct real_add_chain
+{
+    static constexpr int ops_per_cycle {9};
+    static constexpr const char* name {"add"};
+
+    BOOST_SAFE_NUMBERS_HOST_DEVICE static T cycle(T x, const T* y, const T aux)
+    {
+        x = static_cast<T>(x + y[0]);
+        x = static_cast<T>(x + y[1]);
+        x = static_cast<T>(x + y[2]);
+        x = static_cast<T>(x + y[3]);
+        x = static_cast<T>(x + y[4]);
+        x = static_cast<T>(x + y[5]);
+        x = static_cast<T>(x + y[6]);
+        x = static_cast<T>(x + y[7]);
+        x = static_cast<T>(x - aux);
+        return x;
+    }
+};
+
+template <typename T>
+struct real_sub_chain
+{
+    static constexpr int ops_per_cycle {9};
+    static constexpr const char* name {"sub"};
+
+    BOOST_SAFE_NUMBERS_HOST_DEVICE static T cycle(T x, const T* y, const T aux)
+    {
+        x = static_cast<T>(x - y[0]);
+        x = static_cast<T>(x - y[1]);
+        x = static_cast<T>(x - y[2]);
+        x = static_cast<T>(x - y[3]);
+        x = static_cast<T>(x - y[4]);
+        x = static_cast<T>(x - y[5]);
+        x = static_cast<T>(x - y[6]);
+        x = static_cast<T>(x - y[7]);
+        x = static_cast<T>(x + aux);
+        return x;
+    }
+};
+
+// aux is the precomputed product of the eight operands
+template <typename T>
+struct real_mul_chain
+{
+    static constexpr int ops_per_cycle {9};
+    static constexpr const char* name {"mul"};
+
+    BOOST_SAFE_NUMBERS_HOST_DEVICE static T cycle(T x, const T* y, const T aux)
+    {
+        x = static_cast<T>(x * y[0]);
+        x = static_cast<T>(x * y[1]);
+        x = static_cast<T>(x * y[2]);
+        x = static_cast<T>(x * y[3]);
+        x = static_cast<T>(x * y[4]);
+        x = static_cast<T>(x * y[5]);
+        x = static_cast<T>(x * y[6]);
+        x = static_cast<T>(x * y[7]);
+        x = static_cast<T>(x / aux);
+        return x;
+    }
+};
+
+template <typename T>
+struct real_div_chain
+{
+    static constexpr int ops_per_cycle {9};
+    static constexpr const char* name {"div"};
+
+    BOOST_SAFE_NUMBERS_HOST_DEVICE static T cycle(T x, const T* y, const T aux)
+    {
+        x = static_cast<T>(x / y[0]);
+        x = static_cast<T>(x / y[1]);
+        x = static_cast<T>(x / y[2]);
+        x = static_cast<T>(x / y[3]);
+        x = static_cast<T>(x / y[4]);
+        x = static_cast<T>(x / y[5]);
+        x = static_cast<T>(x / y[6]);
+        x = static_cast<T>(x / y[7]);
+        x = static_cast<T>(x * aux);
+        return x;
+    }
+};
+
+// ---------------------------------------------------------------------------
+// Host-side input generation. Every generator fills a builtin vector and a
+// safe vector from the same random draws so both contenders see identical
+// values.
+// ---------------------------------------------------------------------------
+
+inline std::uint64_t draw(std::mt19937_64& rng, const std::uint64_t lo, const std::uint64_t hi)
+{
+    return lo + rng() % (hi - lo + 1U);
+}
+
+inline double draw_real(std::mt19937_64& rng, const double lo, const double hi)
+{
+    return lo + (hi - lo) * (static_cast<double>(rng() >> 11) * 0x1.0p-53);
+}
+
+template <typename TB, typename TS>
+struct paired_vectors
+{
+    std::vector<TB> builtin_vals;
+    std::vector<TS> safe_vals;
+
+    void reserve(const std::size_t n)
+    {
+        builtin_vals.reserve(n);
+        safe_vals.reserve(n);
+    }
+
+    template <typename U>
+    void push(const U v)
+    {
+        builtin_vals.push_back(make_val<TB>(v));
+        safe_vals.push_back(make_val<TS>(v));
+    }
+};
+
+template <typename TB, typename TS>
+struct stream_inputs
+{
+    paired_vectors<TB, TS> a;
+    paired_vectors<TB, TS> b;
+};
+
+// Mirrors the host benchmarks: operands in [1, bits - 1] with a >= b, so
+// add, sub, mul, div, and mod all stay on the non-error path for every width
+template <typename TB, typename TS>
+stream_inputs<TB, TS> gen_stream_int(const int n)
+{
+    const std::uint64_t max_v {sizeof(TB) * 8U - 1U};
+
+    stream_inputs<TB, TS> in;
+    in.a.reserve(static_cast<std::size_t>(n));
+    in.b.reserve(static_cast<std::size_t>(n));
+
+    std::mt19937_64 rng {42};
+    for (int i {0}; i < n; ++i)
+    {
+        auto lhs {draw(rng, 1U, max_v)};
+        auto rhs {draw(rng, 1U, max_v)};
+        if (lhs < rhs)
+        {
+            std::swap(lhs, rhs);
+        }
+        in.a.push(lhs);
+        in.b.push(rhs);
+    }
+
+    return in;
+}
+
+// Mirrors the host float benchmarks: operands in [1, 100]
+template <typename TB, typename TS>
+stream_inputs<TB, TS> gen_stream_real(const int n)
+{
+    stream_inputs<TB, TS> in;
+    in.a.reserve(static_cast<std::size_t>(n));
+    in.b.reserve(static_cast<std::size_t>(n));
+
+    std::mt19937_64 rng {42};
+    for (int i {0}; i < n; ++i)
+    {
+        in.a.push(draw_real(rng, 1.0, 100.0));
+        in.b.push(draw_real(rng, 1.0, 100.0));
+    }
+
+    return in;
+}
+
+template <typename TB, typename TS>
+struct chain_inputs
+{
+    paired_vectors<TB, TS> x0;   // one starting value per thread
+    paired_vectors<TB, TS> ys;   // chain_operands operands per thread
+    paired_vectors<TB, TS> aux;  // one restore value per thread
+};
+
+// Operand ranges follow the int_mixed_chain contract above
+template <typename TB, typename TS>
+chain_inputs<TB, TS> gen_int_chain(const int threads)
+{
+    chain_inputs<TB, TS> in;
+    in.x0.reserve(static_cast<std::size_t>(threads));
+    in.ys.reserve(static_cast<std::size_t>(threads) * chain_operands);
+    in.aux.reserve(static_cast<std::size_t>(threads));
+
+    std::mt19937_64 rng {42};
+    for (int t {0}; t < threads; ++t)
+    {
+        in.x0.push(draw(rng, 0U, 100U));
+        in.ys.push(draw(rng, 11U, 13U));
+        in.ys.push(draw(rng, 8U, 15U));
+        in.ys.push(draw(rng, 1U, 2U));
+        in.ys.push(draw(rng, 1U, 4U));
+        in.ys.push(draw(rng, 1U, 4U));
+        in.ys.push(draw(rng, 5U, 9U));
+        in.ys.push(draw(rng, 1U, 7U));
+        in.ys.push(draw(rng, 1U, 2U));
+        in.aux.push(0U);
+    }
+
+    return in;
+}
+
+enum class real_chain_kind
+{
+    sum,     // operands in [0, 4], aux = sum (add and sub chains)
+    product, // operands in [0.5, 2), aux = product (mul chain)
+    divisor  // operands in [1, 2), aux = product (div chain)
+};
+
+// aux is accumulated in the builtin basis type so the device restore step
+// sees exactly the value the operands were rounded to
+template <typename TB, typename TS>
+chain_inputs<TB, TS> gen_real_chain(const int threads, const real_chain_kind kind)
+{
+    chain_inputs<TB, TS> in;
+    in.x0.reserve(static_cast<std::size_t>(threads));
+    in.ys.reserve(static_cast<std::size_t>(threads) * chain_operands);
+    in.aux.reserve(static_cast<std::size_t>(threads));
+
+    std::mt19937_64 rng {42};
+    for (int t {0}; t < threads; ++t)
+    {
+        in.x0.push(draw_real(rng, 50.0, 100.0));
+
+        auto acc {static_cast<TB>(kind == real_chain_kind::sum ? 0.0 : 1.0)};
+        for (int k {0}; k < chain_operands; ++k)
+        {
+            auto v {0.0};
+            switch (kind)
+            {
+                case real_chain_kind::sum:
+                    v = draw_real(rng, 0.0, 4.0);
+                    break;
+                case real_chain_kind::product:
+                    v = draw_real(rng, 0.5, 2.0);
+                    break;
+                case real_chain_kind::divisor:
+                    v = draw_real(rng, 1.0, 2.0);
+                    break;
+            }
+
+            const auto vb {static_cast<TB>(v)};
+            if (kind == real_chain_kind::sum)
+            {
+                acc = static_cast<TB>(acc + vb);
+            }
+            else
+            {
+                acc = static_cast<TB>(acc * vb);
+            }
+            in.ys.push(vb);
+        }
+
+        in.aux.push(acc);
+    }
+
+    return in;
+}
+
+// ---------------------------------------------------------------------------
+// Verification and reporting
+// ---------------------------------------------------------------------------
+
+template <typename TB, typename TS>
+long long count_mismatches(const std::vector<TB>& builtin_out, const std::vector<TS>& safe_out)
+{
+    long long bad {0};
+    for (std::size_t i {0}; i < builtin_out.size(); ++i)
+    {
+        if (builtin_out[i] != as_builtin<TB>(safe_out[i]))
+        {
+            ++bad;
+        }
+    }
+    return bad;
+}
+
+inline void print_table_header(const char* type_name, const char* shape)
+{
+    std::printf("\n=== %s : %s ===\n", type_name, shape);
+    std::printf("%-8s %16s %16s %8s   %s\n", "op", "builtin (ns/op)", "safe (ns/op)", "ratio", "verify");
+}
+
+// Returns 1 if verification failed so callers can accumulate a failure count
+inline int print_row(const char* op, const double builtin_ns, const double safe_ns, const long long mismatches)
+{
+    const auto ratio {builtin_ns > 0.0 ? safe_ns / builtin_ns : 0.0};
+    if (mismatches == 0)
+    {
+        std::printf("%-8s %16.4f %16.4f %8.2f   ok\n", op, builtin_ns, safe_ns, ratio);
+        return 0;
+    }
+
+    std::printf("%-8s %16.4f %16.4f %8.2f   FAIL (%lld mismatches)\n", op, builtin_ns, safe_ns, ratio, mismatches);
+    return 1;
+}
+
+} // namespace device_bench
+
+#endif // BOOST_SAFE_NUMBERS_BENCHMARKS_DEVICE_BENCH_HPP