From 88a393ec83625b05c4908f6a654a7531a5ba485a Mon Sep 17 00:00:00 2001 From: Matt Borland Date: Fri, 10 Jul 2026 15:15:52 -0400 Subject: [PATCH 1/2] Add initial host vs device benchmarking --- test/benchmarks/CMakeLists.txt | 52 +- test/benchmarks/benchmark_cuda_operations.cu | 431 ++++++++++++++++ test/benchmarks/benchmark_sycl_operations.cpp | 411 +++++++++++++++ test/benchmarks/device_bench.hpp | 475 ++++++++++++++++++ 4 files changed, 1367 insertions(+), 2 deletions(-) create mode 100644 test/benchmarks/benchmark_cuda_operations.cu create mode 100644 test/benchmarks/benchmark_sycl_operations.cpp create mode 100644 test/benchmarks/device_bench.hpp 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 +#include "device_bench.hpp" + +#include + +#include +#include +#include +#include +#include + +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 +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(ptr); + } + + ~device_array() + { + cudaFree(data_); + } + + device_array(const device_array&) = delete; + device_array& operator=(const device_array&) = delete; + + void upload(const std::vector& src) + { + cuda_check(cudaMemcpy(data_, src.data(), src.size() * sizeof(T), cudaMemcpyHostToDevice), "cudaMemcpy to device"); + } + + std::vector download() const + { + std::vector host(n_, device_bench::make_val(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 +__global__ void stream_kernel(const T* a, const T* b, T* out, const int n) +{ + const auto i {static_cast(blockDim.x * blockIdx.x + threadIdx.x)}; + + if (i < n) + { + out[i] = Op::apply(a[i], b[i]); + } +} + +template +__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(blockDim.x * blockIdx.x + threadIdx.x)}; + + if (t < n_threads) + { + const auto base {static_cast(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 +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(ms) * 1.0e6 / launches; +} + +template +int run_stream_op(boost::safe_numbers::device_error_context& ctx, + const device_bench::stream_inputs& host_in, + const device_array& a_b, const device_array& b_b, device_array& out_b, + const device_array& a_s, const device_array& b_s, device_array& 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<<>>(a_b.get(), b_b.get(), out_b.get(), n); + }, cfg.launches)}; + ctx.synchronize(); + + const auto safe_ns {time_launches_ns([&]() + { + stream_kernel<<>>(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(i)], + host_in.b.builtin_vals[static_cast(i)])) + { + ++bad; + } + } + + return device_bench::print_row(Op::name, builtin_ns / n, safe_ns / n, bad); +} + +template