From 0cc3180f26631a98fa3ed53c6d51ef1a617b1ab3 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Thu, 9 Jul 2026 12:25:59 +0200 Subject: [PATCH 01/14] Add pytorch inspired benchmark Fixes: https://github.com/NVIDIA-dev/cccl_private/issues/639 --- cub/benchmarks/CMakeLists.txt | 3 +- .../transform/applications/P1/pytorch.cu | 1262 +++++++++++++++++ 2 files changed, 1264 insertions(+), 1 deletion(-) create mode 100644 cub/benchmarks/bench/transform/applications/P1/pytorch.cu diff --git a/cub/benchmarks/CMakeLists.txt b/cub/benchmarks/CMakeLists.txt index 72fcdd36b50..321c81e26e2 100644 --- a/cub/benchmarks/CMakeLists.txt +++ b/cub/benchmarks/CMakeLists.txt @@ -124,7 +124,8 @@ function(add_bench_dir bench_dir) target_compile_definitions(${base_bench_target} PRIVATE TUNE_BASE=1) target_compile_options( ${base_bench_target} - PRIVATE "$<$:--extended-lambda>" + PRIVATE + "$<$:--extended-lambda;--expt-relaxed-constexpr>" ) if (CUB_ENABLE_TUNING) diff --git a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu new file mode 100644 index 00000000000..9ad515c8e71 --- /dev/null +++ b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu @@ -0,0 +1,1262 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +// NVBench benchmarks for chained elementwise operations put together by Matthias Jouanneaux (DevTech). +// See also: https://github.com/NVIDIA-dev/cccl_private/issues/639 + +// %RANGE% TUNE_BIF_BIAS bif -16:16:4 +// %RANGE% TUNE_ALGORITHM alg 0:4:1 +// %RANGE% TUNE_THREADS tpb 128:1024:128 + +// for TUNE_ALGORITHM == 1 (vectorized), this is the number of vectors per thread, which is similar in spirit +// %RANGE% TUNE_UNROLL_FACTOR unrl 1:4:1 + +// those parameters only apply if TUNE_ALGORITHM == 0 (prefetch) +// %RANGE% TUNE_PREFETCH_MULT pref 1:3:1 + +// those parameters only apply if TUNE_ALGORITHM == 1 (vectorized) +// %RANGE% TUNE_VEC_SIZE_POW2 vsp2 1:6:1 + +#if !TUNE_BASE && TUNE_ALGORITHM != 0 && (TUNE_PREFETCH_MULT != 1) +# error "Non-prefetch algorithms require prefetch multiple to be 1 since they ignore the parameters" +#endif // !TUNE_BASE && TUNE_ALGORITHM != 0 && (TUNE_PREFETCH_MULT != 1) + +#if !TUNE_BASE && TUNE_ALGORITHM != 1 && (TUNE_VEC_SIZE_POW2 != 1) +# error "Non-vectorized algorithms require vector size to be 1 since they ignore the parameters" +#endif // !TUNE_BASE && TUNE_ALGORITHM != 1 && (TUNE_VEC_SIZE_POW2 != 1) + +#include +#include +#include +#include + +#include +#include +#include +#include + +#include "../../common.h" + +// ============================================================================ +// BFloat16 type — replicated from c10::BFloat16 +// (torch/headeronly/util/BFloat16.h) +// ============================================================================ + +namespace bf16_detail +{ +inline __host__ __device__ float f32_from_bits(uint16_t src) +{ + float res = 0; + uint32_t tmp = src; + tmp <<= 16; + std::memcpy(&res, &tmp, sizeof(tmp)); + return res; +} + +inline __host__ __device__ uint16_t round_to_nearest_even(float src) +{ + if (std::isnan(src)) + { + return UINT16_C(0x7FC0); + } + else + { + uint32_t U32; + std::memcpy(&U32, &src, sizeof(U32)); + uint32_t rounding_bias = ((U32 >> 16) & 1) + UINT32_C(0x7FFF); + return static_cast((U32 + rounding_bias) >> 16); + } +} +} // namespace bf16_detail + +struct alignas(2) BFloat16 +{ + uint16_t x; + + BFloat16() = default; + + struct from_bits_t + {}; + static constexpr __host__ __device__ from_bits_t from_bits() + { + return from_bits_t(); + } + + constexpr __host__ __device__ BFloat16(unsigned short bits, from_bits_t) + : x(bits) + {} + + /* implicit */ inline __host__ __device__ BFloat16(float value); + inline __host__ __device__ operator float() const; + +#if defined(__CUDACC__) + inline __host__ __device__ BFloat16(const __nv_bfloat16& value); + explicit inline __host__ __device__ operator __nv_bfloat16() const; +#endif +}; + +inline __host__ __device__ BFloat16::BFloat16(float value) + : +#if defined(__CUDACC__) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800 + x(__bfloat16_as_ushort(__float2bfloat16(value))) +#else + x(bf16_detail::round_to_nearest_even(value)) +#endif +{} + +inline __host__ __device__ BFloat16::operator float() const +{ +#if defined(__CUDACC__) + return __bfloat162float(*reinterpret_cast(&x)); +#else + return bf16_detail::f32_from_bits(x); +#endif +} + +#if defined(__CUDACC__) +inline __host__ __device__ BFloat16::BFloat16(const __nv_bfloat16& value) +{ + x = *reinterpret_cast(&value); +} +inline __host__ __device__ BFloat16::operator __nv_bfloat16() const +{ + return *reinterpret_cast(&x); +} +#endif + +// Arithmetic — BFloat16 x BFloat16 → BFloat16 +inline __host__ __device__ BFloat16 operator+(const BFloat16& a, const BFloat16& b) +{ + return static_cast(a) + static_cast(b); +} +inline __host__ __device__ BFloat16 operator-(const BFloat16& a, const BFloat16& b) +{ + return static_cast(a) - static_cast(b); +} +inline __host__ __device__ BFloat16 operator*(const BFloat16& a, const BFloat16& b) +{ + return static_cast(a) * static_cast(b); +} +inline __host__ __device__ BFloat16 operator/(const BFloat16& a, const BFloat16& b) +{ + return static_cast(a) / static_cast(b); +} +inline __host__ __device__ BFloat16 operator-(const BFloat16& a) +{ + return -static_cast(a); +} + +// Compound assignment — BFloat16 +inline __host__ __device__ BFloat16& operator+=(BFloat16& a, const BFloat16& b) +{ + a = a + b; + return a; +} +inline __host__ __device__ BFloat16& operator-=(BFloat16& a, const BFloat16& b) +{ + a = a - b; + return a; +} +inline __host__ __device__ BFloat16& operator*=(BFloat16& a, const BFloat16& b) +{ + a = a * b; + return a; +} +inline __host__ __device__ BFloat16& operator/=(BFloat16& a, const BFloat16& b) +{ + a = a / b; + return a; +} + +// Arithmetic — BFloat16 x float → float +inline __host__ __device__ float operator+(BFloat16 a, float b) +{ + return static_cast(a) + b; +} +inline __host__ __device__ float operator-(BFloat16 a, float b) +{ + return static_cast(a) - b; +} +inline __host__ __device__ float operator*(BFloat16 a, float b) +{ + return static_cast(a) * b; +} +inline __host__ __device__ float operator/(BFloat16 a, float b) +{ + return static_cast(a) / b; +} +inline __host__ __device__ float operator+(float a, BFloat16 b) +{ + return a + static_cast(b); +} +inline __host__ __device__ float operator-(float a, BFloat16 b) +{ + return a - static_cast(b); +} +inline __host__ __device__ float operator*(float a, BFloat16 b) +{ + return a * static_cast(b); +} +inline __host__ __device__ float operator/(float a, BFloat16 b) +{ + return a / static_cast(b); +} + +// Compound assignment — float x BFloat16 → float +inline __host__ __device__ float& operator+=(float& a, const BFloat16& b) +{ + return a += static_cast(b); +} +inline __host__ __device__ float& operator-=(float& a, const BFloat16& b) +{ + return a -= static_cast(b); +} +inline __host__ __device__ float& operator*=(float& a, const BFloat16& b) +{ + return a *= static_cast(b); +} +inline __host__ __device__ float& operator/=(float& a, const BFloat16& b) +{ + return a /= static_cast(b); +} + +// Arithmetic — BFloat16 x int → BFloat16 +inline __host__ __device__ BFloat16 operator+(BFloat16 a, int b) +{ + return a + static_cast(static_cast(b)); +} +inline __host__ __device__ BFloat16 operator-(BFloat16 a, int b) +{ + return a - static_cast(static_cast(b)); +} +inline __host__ __device__ BFloat16 operator*(BFloat16 a, int b) +{ + return a * static_cast(static_cast(b)); +} +inline __host__ __device__ BFloat16 operator/(BFloat16 a, int b) +{ + return a / static_cast(static_cast(b)); +} +inline __host__ __device__ BFloat16 operator+(int a, BFloat16 b) +{ + return static_cast(static_cast(a)) + b; +} +inline __host__ __device__ BFloat16 operator-(int a, BFloat16 b) +{ + return static_cast(static_cast(a)) - b; +} +inline __host__ __device__ BFloat16 operator*(int a, BFloat16 b) +{ + return static_cast(static_cast(a)) * b; +} +inline __host__ __device__ BFloat16 operator/(int a, BFloat16 b) +{ + return static_cast(static_cast(a)) / b; +} + +// Comparison — for std::min/std::max +inline __host__ __device__ bool operator>(BFloat16& lhs, BFloat16& rhs) +{ + return float(lhs) > float(rhs); +} +inline __host__ __device__ bool operator<(BFloat16& lhs, BFloat16& rhs) +{ + return float(lhs) < float(rhs); +} + +// NVBench type registration +NVBENCH_DECLARE_TYPE_STRINGS(BFloat16, "bf16", "BFloat16"); + +// ============================================================================ +// at::opmath_type — the compute type for intermediate math +// float for both float and BFloat16 (ATen/OpMathType.h) +// ============================================================================ + +template +struct opmath_type_impl +{ + using type = T; +}; +template <> +struct opmath_type_impl +{ + using type = float; +}; +template +using opmath_type = typename opmath_type_impl::type; + +// ============================================================================ +// Replicate ATen/c10 helpers without ATen dependencies. +// Each wrapper is annotated with the ATen source it replicates. +// ============================================================================ + +// at::_isnan (ATen/NumericUtils.h:30) — on CUDA this is just ::isnan +template +__device__ __forceinline__ bool _isnan(T val) +{ + return ::isnan(static_cast(val)); +} + +// c10::cuda::compat (c10/cuda/CUDAMathCompat.h) +namespace compat +{ +__device__ __forceinline__ float copysign(float x, float y) +{ + return ::copysignf(x, y); +} +__device__ __forceinline__ double copysign(double x, double y) +{ + return ::copysign(x, y); +} +__device__ __forceinline__ float exp(float x) +{ + return ::expf(x); +} +__device__ __forceinline__ double exp(double x) +{ + return ::exp(x); +} +__device__ __forceinline__ float log1p(float x) +{ + return ::log1pf(x); +} +__device__ __forceinline__ double log1p(double x) +{ + return ::log1p(x); +} +__device__ __forceinline__ float tanh(float x) +{ + return ::tanhf(x); +} +__device__ __forceinline__ double tanh(double x) +{ + return ::tanh(x); +} +} // namespace compat + +// c10::div_floor_floating (c10/util/generic_math.h:34) +template +__device__ __forceinline__ scalar_t div_floor_floating(scalar_t a, scalar_t b) +{ + if (b == 0) + { + return a / b; + } + + auto mod = std::fmod(a, b); + auto div = (a - mod) / b; + if ((mod != 0) && (b < 0) != (mod < 0)) + { + div -= scalar_t(1); + } + + scalar_t floordiv; + if (div != 0) + { + floordiv = std::floor(div); + if (div - floordiv > scalar_t(0.5)) + { + floordiv += scalar_t(1.0); + } + } + else + { + floordiv = compat::copysign(scalar_t(0), a / b); + } + return floordiv; +} + +// pow_ (native/cuda/Pow.cuh:40) +template +__device__ __forceinline__ Base_type pow_(Base_type base, Exp_type exp) +{ + return ::pow(base, exp); +} + +// is_lerp_weight_small + lerp (native/Lerp.h:11,21) +template +__device__ __forceinline__ bool is_lerp_weight_small(scalar_t weight) +{ + return std::abs(weight) < scalar_t(0.5); +} + +template +__device__ __forceinline__ scalar_t lerp(scalar_t self_, scalar_t end_, weight_t weight_) +{ + using opmath_t = opmath_type; + using opmath_weight_t = opmath_type; + + opmath_t self = self_; + opmath_t end = end_; + opmath_weight_t weight = weight_; + + return is_lerp_weight_small(weight) ? self + weight * (end - self) : end - (end - self) * (opmath_t(1) - weight); +} + +// pointwise_op_impl (native/cuda/DeviceAddCmulCdiv.cuh:9) +template +__device__ __forceinline__ opmath_t +pointwise_op_impl(opmath_t input, opmath_t tensor1, opmath_t tensor2, opmath_t alpha, Op op) +{ + if (alpha == opmath_t(1)) + { + if constexpr (std::is_same_v> && std::is_floating_point_v) + { + return std::fma(tensor1, tensor2, input); + } + else + { + return input + op(tensor1, tensor2); + } + } + if constexpr (std::is_floating_point_v) + { + return std::fma(alpha, op(tensor1, tensor2), input); + } + else + { + return input + alpha * op(tensor1, tensor2); + } +} + +// DivFunctor (native/cuda/BinaryInternal.h:20) +template +struct DivFunctor +{ + __device__ scalar_t operator()(scalar_t a, scalar_t b) const + { + return a / b; + } +}; + +// MulFunctor (native/cuda/BinaryInternal.h:27) +template +struct MulFunctor +{ + __device__ T operator()(T a, T b) const + { + return a * b; + } +}; + +// CUDAFunctorOnSelf_add — torchgen-generated ufunc functor for add(tensor, scalar) +// (torchgen/dest/ufunc.py, native/ufunc/add.h:14) +template +struct CUDAFunctorOnSelf_add +{ + using opmath_t = opmath_type; + opmath_t other_; + opmath_t alpha_; + CUDAFunctorOnSelf_add(opmath_t other, opmath_t alpha) + : other_(other) + , alpha_(alpha) + {} + __device__ scalar_t operator()(scalar_t self) const + { + return static_cast(self) + alpha_ * other_; + } +}; + +// CUDAFunctor_add — torchgen-generated ufunc functor for add(tensor, tensor) +// (torchgen/dest/ufunc.py, native/ufunc/add.h:14) +template +struct CUDAFunctor_add +{ + using opmath_t = opmath_type; + opmath_t alpha_; + CUDAFunctor_add(opmath_t alpha) + : alpha_(alpha) + {} + __device__ scalar_t operator()(scalar_t self, scalar_t other) const + { + return static_cast(self) + alpha_ * static_cast(other); + } +}; + +// AbsFunctor (native/cuda/AbsKernel.cu:11) +template +struct AbsFunctor +{ + __device__ __forceinline__ scalar_t operator()(const scalar_t a) const + { + return std::abs(a); + } +}; + +// CompareFunctor (native/cuda/CompareKernels.cu:14 / 17) +enum class OpType +{ + GE, + GT, + LE, + LT +}; + +template +struct CompareFunctor +{ + constexpr CompareFunctor(OpType op) + : op_(op) {}; + OpType op_; + __device__ __forceinline__ bool operator()(scalar_t a, scalar_t b) const + { + if (op_ == OpType::GE) + { + return a >= b; + } + else if (op_ == OpType::GT) + { + return a > b; + } + else if (op_ == OpType::LE) + { + return a <= b; + } + else + { // LT + return a < b; + } + } +}; + +// ============================================================================ +// RNG helpers +// ============================================================================ + +template +struct normal_gen +{ + float mean, stddev; + int offset; + __host__ __device__ T operator()(int idx) const + { + thrust::default_random_engine rng; + thrust::random::normal_distribution dist(mean, stddev); + rng.discard(offset + idx); + return T(dist(rng)); + } +}; + +template +void fill_normal(thrust::device_vector& v, OffsetT n, int buf_idx) +{ + v.resize(n); + thrust::transform(thrust::counting_iterator(0), + thrust::counting_iterator(n), + v.begin(), + normal_gen{0.0f, 1.0f, buf_idx * static_cast(n)}); +} + +// ============================================================================ +// Helper to call DeviceTransform::Transform with the tuning policy +// ============================================================================ + +template +void transform(::cuda::std::tuple inputs, Output output, int n, TransformOp op, cudaStream_t stream) +{ + auto env = cuda::std::execution::env{ + ::cuda::stream_ref{stream} +#if !TUNE_BASE + , + cuda::execution::tune(policy_selector{}) +#endif // !TUNE_BASE + }; + cub::DeviceTransform::Transform(inputs, output, n, op, env); +} + +template +void transform(Input input, Output output, int n, TransformOp op, cudaStream_t stream) +{ + transform(::cuda::std::make_tuple(input), output, n, op, stream); +} + +// ============================================================================ +// Element types +// ============================================================================ + +#ifdef TUNE_T +using element_types = nvbench::type_list; +#else +using element_types = nvbench::type_list; +#endif + +// note: pytorch always uses int for offset type +// but be careful not to use too large of a range for the Elements axis, +// otherwise, we will run into overflows. +using pytorch_offset_types = nvbench::type_list; + +// ============================================================================ +// chained_eltwise_many_in_many_inst +// +// div_floor -> div_trunc -> div -> atan2 -> hypot -> +// xlogy -> xlog1py -> logaddexp -> logaddexp2 -> pow +// 11 inputs, 10 binary ops +// ============================================================================ + +template +static void chained_eltwise_many_in_many_inst(nvbench::state& state, nvbench::type_list) +try +{ + const auto n = static_cast(state.get_int64("Elements{io}")); + + constexpr int num_in = 11; + thrust::device_vector in[num_in]; + for (int i = 0; i < num_in; i++) + { + fill_normal(in[i], n, i); + } + thrust::device_vector tmpA(n), tmpB(n); + + T* d_in[num_in]; + for (int i = 0; i < num_in; i++) + { + d_in[i] = thrust::raw_pointer_cast(in[i].data()); + } + T* d_a = thrust::raw_pointer_cast(tmpA.data()); + T* d_b = thrust::raw_pointer_cast(tmpB.data()); + + state.add_element_count(n); + state.add_global_memory_reads(20L * int64_t{n}); + state.add_global_memory_writes(10L * int64_t{n}); + + // logaddexp2 captures inv_log_2 — native/cuda/LogAddExpKernel.cu:272 + using opmath_t = opmath_type; + const auto inv_log_2 = static_cast(1.0 / 0.693147180559945309417232121458176); + + state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch, [&](const nvbench::launch& launch) { + auto s = launch.get_stream().get_stream(); + + // div_floor: native/cuda/BinaryDivFloorKernel.cu:72, helper c10/util/generic_math.h:34 + transform( + ::cuda::std::make_tuple(d_in[0], d_in[1]), + d_a, + n, + [] __device__(T a, T b) -> T { + return div_floor_floating(a, b); + }, + s); + + // div_trunc: native/cuda/BinaryDivTruncKernel.cu:42 + transform( + ::cuda::std::make_tuple(d_a, d_in[2]), + d_b, + n, + [] __device__(T a, T b) -> T { + return std::trunc(a / b); + }, + s); + + // div: native/cuda/BinaryDivTrueKernel.cu:54, DivFunctor in native/cuda/BinaryInternal.h:20 + transform(::cuda::std::make_tuple(d_b, d_in[3]), d_a, n, DivFunctor(), s); + + // atan2: native/cuda/BinaryGeometricKernels.cu:18 + transform( + ::cuda::std::make_tuple(d_a, d_in[4]), + d_b, + n, + [] __device__(T a, T b) -> T { + return ::atan2(a, b); + }, + s); + + // hypot: native/cuda/BinaryGeometricKernels.cu:29 + transform( + ::cuda::std::make_tuple(d_b, d_in[5]), + d_a, + n, + [] __device__(T a, T b) -> T { + return ::hypot(a, b); + }, + s); + + // xlogy: native/cuda/BinaryMiscOpsKernels.cu:46 + transform( + ::cuda::std::make_tuple(d_a, d_in[6]), + d_b, + n, + [] __device__(T x, T y) -> T { + if (_isnan(y)) + { + return NAN; + } + if (x == 0) + { + return 0; + } + return x * std::log(y); + }, + s); + + // xlog1py: native/cuda/BinaryMiscOpsKernels.cu:60 + transform( + ::cuda::std::make_tuple(d_b, d_in[7]), + d_a, + n, + [] __device__(T x, T y) -> T { + if (_isnan(y)) + { + return NAN; + } + if (x == 0) + { + return 0; + } + return x * std::log1p(y); + }, + s); + + // logaddexp: native/cuda/LogAddExpKernel.cu:253 + transform( + ::cuda::std::make_tuple(d_a, d_in[8]), + d_b, + n, + [] __device__(T a_, T b_) -> T { + using opmath_t = opmath_type; + const auto a = static_cast(a_); + const auto b = static_cast(b_); + if (::isinf(a) && a == b) + { + return a; + } + else + { + const auto m = ::max(a, b); + return m + ::log1p(::exp(-::abs(a - b))); + } + }, + s); + + // logaddexp2: native/cuda/LogAddExpKernel.cu:272 + transform( + ::cuda::std::make_tuple(d_b, d_in[9]), + d_a, + n, + [inv_log_2] __device__(T a_, T b_) -> T { + using opmath_t = opmath_type; + const auto a = static_cast(a_); + const auto b = static_cast(b_); + if (::isinf(a) && a == b) + { + return a; + } + else + { + const auto m = ::max(a, b); + return m + ::log1p(::exp2(-::abs(a - b))) * inv_log_2; + } + }, + s); + + // pow (tensor,tensor): native/cuda/PowKernel.cu:136, helper native/cuda/Pow.cuh:40 + transform( + ::cuda::std::make_tuple(d_a, d_in[10]), + d_b, + n, + [] __device__(T base, T exp) -> T { + return pow_(base, exp); + }, + s); + }); +} +catch (const std::bad_alloc&) +{ + state.skip("Skipping: out of memory."); +} + +// ============================================================================ +// chained_eltwise_many_in_few_inst +// +// mse_loss -> smooth_l1_loss -> huber_loss -> clamp_min -> +// mul -> add -> addcmul -> lerp(scalar) -> lerp(tensor) -> greater +// 13 inputs, 8 binary ops + 2 ternary ops +// ============================================================================ + +template +static void chained_eltwise_many_in_few_inst(nvbench::state& state, nvbench::type_list) +try +{ + const auto n = static_cast(state.get_int64("Elements{io}")); + + constexpr int num_in = 13; + thrust::device_vector in[num_in]; + for (int i = 0; i < num_in; i++) + { + fill_normal(in[i], n, i); + } + thrust::device_vector tmpA(n), tmpB(n); + + T* d_in[num_in]; + for (int i = 0; i < num_in; i++) + { + d_in[i] = thrust::raw_pointer_cast(in[i].data()); + } + T* d_a = thrust::raw_pointer_cast(tmpA.data()); + T* d_b = thrust::raw_pointer_cast(tmpB.data()); + + // 8 binary (16 reads) + 2 ternary (6 reads) = 22 reads, 10 writes + state.add_element_count(n); + state.add_global_memory_reads(22L * int64_t{n}); + state.add_global_memory_writes(10L * int64_t{n}); + + // Captured scalar parameters, matching how ATen sets them up before gpu_kernel + using opmath_t = opmath_type; + T beta_val(1.0); // smooth_l1: scalar_t beta_val(beta) + T delta_val(1.0); // huber: scalar_t delta_val(delta) + // note: opmath_type is same as at::acc_type here + using accscalar_t = opmath_type; // addcmul: at::acc_type + auto alpha = accscalar_t(1); // addcmul: value.to() + auto weight_val = opmath_t(4.0); // lerp scalar: weight.to() + + state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch, [&](const nvbench::launch& launch) { + auto s = launch.get_stream().get_stream(); + + // mse_loss: native/cuda/BinaryMiscOpsKernels.cu:37 + transform( + ::cuda::std::make_tuple(d_in[0], d_in[1]), + d_a, + n, + [] __device__(T a, T b) -> T { + auto diff = a - b; + return diff * diff; + }, + s); + + // smooth_l1_loss(beta=1.0): native/cuda/BinaryMiscOpsKernels.cu:19 + transform( + ::cuda::std::make_tuple(d_a, d_in[2]), + d_b, + n, + [beta_val] __device__(T a, T b) -> T { + auto z = ::abs(a - b); + return z < beta_val ? T(0.5) * z * z / beta_val : z - T(0.5) * beta_val; + }, + s); + + // huber_loss(delta=1.0): native/cuda/BinaryMiscOpsKernels.cu:29 + transform( + ::cuda::std::make_tuple(d_b, d_in[3]), + d_a, + n, + [delta_val] __device__(T a, T b) -> T { + auto z = ::abs(a - b); + return z < delta_val ? T(0.5) * z * z : delta_val * (z - T(0.5) * delta_val); + }, + s); + + // clamp(min=tensor) -> maximum: native/cuda/MaxMinElementwiseKernel.cu:28 + transform( + ::cuda::std::make_tuple(d_a, d_in[4]), + d_b, + n, + [] __device__(T a, T b) -> T { + if (a != a) + { + return a; + } + else if (b != b) + { + return b; + } + else + { + return ::max(a, b); + } + }, + s); + + // mul: native/cuda/BinaryMulKernel.cu:39, MulFunctor in native/cuda/BinaryInternal.h:27 + using mul_opmath_t = opmath_type; + transform(::cuda::std::make_tuple(d_b, d_in[5]), d_a, n, MulFunctor(), s); + + // add(alpha=1): native/ufunc/add.h:14, torchgen/dest/ufunc.py + transform(::cuda::std::make_tuple(d_a, d_in[6]), d_b, n, CUDAFunctor_add(1.0), s); + + // addcmul(value=1): native/cuda/PointwiseOpsKernel.cu:87, native/cuda/DeviceAddCmulCdiv.cuh:9 + transform( + ::cuda::std::make_tuple(d_b, d_in[7], d_in[8]), + d_a, + n, + [alpha] __device__(T a, T b, T c) -> T { + return pointwise_op_impl(a, b, c, alpha, std::multiplies()); + }, + s); + + // lerp(weight=4.0): native/cuda/Lerp.cu:130, native/Lerp.h:21 + transform( + ::cuda::std::make_tuple(d_a, d_in[9]), + d_b, + n, + [=] __device__(T self_val, T end_val) { + return lerp(self_val, end_val, weight_val); + }, + s); + + // lerp(weight=tensor): native/cuda/Lerp.cu:76, native/Lerp.h:21 + transform( + ::cuda::std::make_tuple(d_b, d_in[10], d_in[11]), + d_a, + n, + [] __device__(T self_val, T end_val, T weight_val) -> T { + return lerp(self_val, end_val, weight_val); + }, + s); + + // note: even though output is bool, we use d_b as output because + // it must hold at least enough memory per element for bool (1 byte) + // greater: native/cuda/CompareKernels.cu:69 + CompareFunctor comp_f(OpType::GT); + transform(::cuda::std::make_tuple(d_a, d_in[12]), d_b, n, comp_f, s); + }); +} +catch (const std::bad_alloc&) +{ + state.skip("Skipping: out of memory."); +} + +// ============================================================================ +// chained_eltwise_few_in_many_inst +// +// pow(2.5) -> tanh -> sin -> cos -> softplus -> +// silu -> mish -> elu -> gelu -> logsigmoid +// 1 input, 10 unary ops +// ============================================================================ + +template +static void chained_eltwise_few_in_many_inst(nvbench::state& state, nvbench::type_list) +try +{ + const auto n = static_cast(state.get_int64("Elements{io}")); + + thrust::device_vector input(n); + fill_normal(input, n, 0); + thrust::device_vector tmpA(n), tmpB(n); + + T* d_in = thrust::raw_pointer_cast(input.data()); + T* d_a = thrust::raw_pointer_cast(tmpA.data()); + T* d_b = thrust::raw_pointer_cast(tmpB.data()); + + state.add_element_count(n); + state.add_global_memory_reads(10L * int64_t{n}); + state.add_global_memory_writes(10L * int64_t{n}); + + // Captured scalar parameters + using opmath_t = opmath_type; + const auto exp_val = T(2.5); // pow: exp_scalar.to() + auto beta = opmath_t(1); // softplus: beta_.to() + auto threshold = opmath_t(20); // softplus: threshold_.to() + auto negcoef = opmath_t(1) * opmath_t(1); // elu: alpha * scale + auto poscoef = opmath_t(1); // elu: scale + auto negiptcoef = opmath_t(1); // elu: input_scale + + state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch, [&](const nvbench::launch& launch) { + auto s = launch.get_stream().get_stream(); + + // pow(scalar=2.5): native/cuda/PowKernel.cu:163, helper native/cuda/Pow.cuh:40 + transform( + d_in, + d_a, + n, + [=] __device__(T base) -> T { + return pow_(base, exp_val); + }, + s); + + // tanh: native/cuda/UnaryGeometricTanhKernel.cu:50 + transform( + d_a, + d_b, + n, + [] __device__(T a) -> T { + return ::tanh(a); + }, + s); + + // sin: native/cuda/UnaryGeometricSinKernel.cu:50 + transform( + d_b, + d_a, + n, + [] __device__(T a) -> T { + return ::sin(a); + }, + s); + + // cos: native/cuda/UnaryGeometricCosKernel.cu:50 + transform( + d_a, + d_b, + n, + [] __device__(T a) -> T { + return ::cos(a); + }, + s); + + // softplus(beta=1, threshold=20): native/cuda/ActivationSoftplusKernel.cu:35 + transform( + d_b, + d_a, + n, + [beta, threshold] __device__(T a) -> T { + using opmath_t = opmath_type; + opmath_t aop = static_cast(a); + return (aop * beta) > threshold ? aop : (::log1p(std::exp(aop * beta))) / beta; + }, + s); + + // silu: native/cuda/ActivationSiluKernel.cu:30 + transform( + d_a, + d_b, + n, + [] __device__(T x) -> T { + using opmath_t = opmath_type; + const opmath_t x_acc = static_cast(x); + return x_acc / (opmath_t(1) + ::exp(-x_acc)); + }, + s); + + // mish: native/cuda/ActivationMishKernel.cu:29 + transform( + d_b, + d_a, + n, + [] __device__(T x) -> T { + using opmath_t = opmath_type; + const opmath_t x_acc = static_cast(x); + return x_acc * compat::tanh(compat::log1p(compat::exp(x_acc))); + }, + s); + + // elu(alpha=1, scale=1, input_scale=1): native/cuda/ActivationEluKernel.cu:37 + transform( + d_a, + d_b, + n, + [negcoef, poscoef, negiptcoef] __device__(T a) -> T { + using opmath_t = opmath_type; + opmath_t aop = static_cast(a); + return aop > 0 ? aop * poscoef : std::expm1(aop * negiptcoef) * negcoef; + }, + s); + + // gelu(approximate='none'): native/cuda/ActivationGeluKernel.cu:35 + transform( + d_b, + d_a, + n, + [] __device__(T x) -> T { + using opmath_t = opmath_type; + constexpr opmath_t kAlpha = M_SQRT1_2; + return static_cast(x) * opmath_t(0.5) * (opmath_t(1) + ::erf(static_cast(x) * kAlpha)); + }, + s); + + // logsigmoid: native/cuda/ActivationLogSigmoidKernel.cu:30 + transform( + d_a, + d_b, + n, + [] __device__(T in_) -> T { + using opmath_t = opmath_type; + const opmath_t in = in_; + const auto min = std::min(opmath_t(0), in); + const auto z = std::exp(-std::abs(in)); + return min - std::log1p(z); + }, + s); + }); +} +catch (const std::bad_alloc&) +{ + state.skip("Skipping: out of memory."); +} + +// ============================================================================ +// chained_eltwise_few_in_few_inst +// +// add(0.5) -> neg -> clamp(-2,1) -> abs -> mul(1.5) -> +// leaky_relu -> hardswish -> hardshrink -> hardsigmoid -> gt(0) +// 1 input, 10 unary ops +// ============================================================================ + +template +static void chained_eltwise_few_in_few_inst(nvbench::state& state, nvbench::type_list) +try +{ + const auto n = static_cast(state.get_int64("Elements{io}")); + + thrust::device_vector input(n); + fill_normal(input, n, 0); + thrust::device_vector tmpA(n), tmpB(n); + + T* d_in = thrust::raw_pointer_cast(input.data()); + T* d_a = thrust::raw_pointer_cast(tmpA.data()); + T* d_b = thrust::raw_pointer_cast(tmpB.data()); + + state.add_element_count(n); + state.add_global_memory_reads(10L * int64_t{n}); + state.add_global_memory_writes(10L * int64_t{n}); + + // Captured scalar parameters + using opmath_t = opmath_type; + + // clamp: native/cuda/TensorCompare.cu:58 + const auto lim0_val = opmath_t(-2); + const auto lim1_val = opmath_t(1); + const auto minmax = 2; // 0=Min, 1=Max, 2=MinMax + + // mul scalar: MulFunctor via BUnaryFunctor with captured scalar + const auto mul_scalar = opmath_t(1.5); + + // leaky_relu: native/cuda/ActivationLeakyReluKernel.cu:31 + auto negval = opmath_t(0.01); // negval_.to() + + // hardswish: native/cuda/ActivationHardswishKernel.cu:25 + const opmath_t zero(0.0f); + const opmath_t one_sixth(1.0f / 6.0f); + const opmath_t three(3.0f); + const opmath_t six(6.0f); + + // hardshrink: native/cuda/ActivationHardshrinkKernel.cu:29 + auto lambd = T(0.5); // value.to() + + // gt scalar: native/cuda/CompareKernels.cu:47 + const T rhs(0); + + state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch, [&](const nvbench::launch& launch) { + auto s = launch.get_stream().get_stream(); + + // add(scalar, alpha=1): native/ufunc/add.h:14 + transform(d_in, d_a, n, CUDAFunctorOnSelf_add(T(0.5), T(1)), s); + + // neg: native/cuda/UnarySignKernels.cu:54 + transform( + d_a, + d_b, + n, + [] __device__(T a) -> T { + return -a; + }, + s); + + // clamp(min=-2, max=1): native/cuda/TensorCompare.cu:58 (MinMax branch) + transform( + d_b, + d_a, + n, + [=] __device__(T v) -> T { + using opmath_t = opmath_type; + if (_isnan(static_cast(v))) + { + return v; + } + else if (minmax == 0) + { + return ::max(static_cast(v), lim0_val); + } + else if (minmax == 1) + { + return ::min(static_cast(v), lim0_val); + } + else + { + return ::min(::max(static_cast(v), lim0_val), lim1_val); + } + }, + s); + + // abs: native/cuda/AbsKernel.cu:39, AbsFunctor:11 + transform(d_a, d_b, n, AbsFunctor(), s); + + // mul(scalar=1.5): native/cuda/BinaryMulKernel.cu:39, MulFunctor via BUnaryFunctor + transform( + d_b, + d_a, + n, + [mul_scalar] __device__(T a) -> T { + return MulFunctor()(a, mul_scalar); + }, + s); + + // leaky_relu(slope=0.01): native/cuda/ActivationLeakyReluKernel.cu:31 + transform( + d_a, + d_b, + n, + [negval] __device__(T a) -> T { + using opmath_t = opmath_type; + opmath_t aop = static_cast(a); + return aop > opmath_t(0) ? aop : aop * negval; + }, + s); + + // hardswish: native/cuda/ActivationHardswishKernel.cu:25 + transform( + d_b, + d_a, + n, + [zero, one_sixth, three, six] __device__(T self_val) -> T { + using opmath_t = opmath_type; + opmath_t x = static_cast(self_val); + return x * std::min(std::max(x + three, zero), six) * one_sixth; + }, + s); + + // hardshrink(lambd=0.5): native/cuda/ActivationHardshrinkKernel.cu:29 + transform( + d_a, + d_b, + n, + [lambd] __device__(T a) -> T { + return (a >= -lambd && a <= lambd) ? T(0) : a; + }, + s); + + // hardsigmoid: native/cuda/ActivationHardsigmoidKernel.cu:30 + transform( + d_b, + d_a, + n, + [zero, one_sixth, three, six] __device__(T self_val) -> T { + using opmath_t = opmath_type; + opmath_t x = static_cast(self_val); + return std::min(std::max(x + three, zero), six) * one_sixth; + }, + s); + + // gt(scalar=0): native/cuda/CompareKernels.cu:47 + CompareFunctor comp_f(OpType::GT); + transform( + d_a, + d_b, + n, + [=] __device__(T lhs) -> T { + return comp_f(lhs, rhs); + }, + s); + }); +} +catch (const std::bad_alloc&) +{ + state.skip("Skipping: out of memory."); +} + +NVBENCH_BENCH_TYPES(chained_eltwise_many_in_many_inst, NVBENCH_TYPE_AXES(element_types, pytorch_offset_types)) + .set_name("chained_eltwise_many_in_many_inst") + .set_type_axes_names({"T{ct}", "OffsetT{ct}"}) + .add_int64_power_of_two_axis("Elements{io}", nvbench::range(16, 28, 4)); + +NVBENCH_BENCH_TYPES(chained_eltwise_many_in_few_inst, NVBENCH_TYPE_AXES(element_types, pytorch_offset_types)) + .set_name("chained_eltwise_many_in_few_inst") + .set_type_axes_names({"T{ct}", "OffsetT{ct}"}) + .add_int64_power_of_two_axis("Elements{io}", nvbench::range(16, 28, 4)); + +NVBENCH_BENCH_TYPES(chained_eltwise_few_in_many_inst, NVBENCH_TYPE_AXES(element_types, pytorch_offset_types)) + .set_name("chained_eltwise_few_in_many_inst") + .set_type_axes_names({"T{ct}", "OffsetT{ct}"}) + .add_int64_power_of_two_axis("Elements{io}", nvbench::range(16, 28, 4)); + +NVBENCH_BENCH_TYPES(chained_eltwise_few_in_few_inst, NVBENCH_TYPE_AXES(element_types, pytorch_offset_types)) + .set_name("chained_eltwise_few_in_few_inst") + .set_type_axes_names({"T{ct}", "OffsetT{ct}"}) + .add_int64_power_of_two_axis("Elements{io}", nvbench::range(16, 28, 4)); From cc0b391504e6338e098f8658a2bf1f7f3cf7a785 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Thu, 9 Jul 2026 13:01:00 +0200 Subject: [PATCH 02/14] Drop compat math functions --- .../transform/applications/P1/pytorch.cu | 41 +------------------ 1 file changed, 2 insertions(+), 39 deletions(-) diff --git a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu index 9ad515c8e71..d44cfb58b80 100644 --- a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu +++ b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu @@ -297,43 +297,6 @@ __device__ __forceinline__ bool _isnan(T val) return ::isnan(static_cast(val)); } -// c10::cuda::compat (c10/cuda/CUDAMathCompat.h) -namespace compat -{ -__device__ __forceinline__ float copysign(float x, float y) -{ - return ::copysignf(x, y); -} -__device__ __forceinline__ double copysign(double x, double y) -{ - return ::copysign(x, y); -} -__device__ __forceinline__ float exp(float x) -{ - return ::expf(x); -} -__device__ __forceinline__ double exp(double x) -{ - return ::exp(x); -} -__device__ __forceinline__ float log1p(float x) -{ - return ::log1pf(x); -} -__device__ __forceinline__ double log1p(double x) -{ - return ::log1p(x); -} -__device__ __forceinline__ float tanh(float x) -{ - return ::tanhf(x); -} -__device__ __forceinline__ double tanh(double x) -{ - return ::tanh(x); -} -} // namespace compat - // c10::div_floor_floating (c10/util/generic_math.h:34) template __device__ __forceinline__ scalar_t div_floor_floating(scalar_t a, scalar_t b) @@ -361,7 +324,7 @@ __device__ __forceinline__ scalar_t div_floor_floating(scalar_t a, scalar_t b) } else { - floordiv = compat::copysign(scalar_t(0), a / b); + floordiv = ::copysignf(scalar_t(0), a / b); } return floordiv; } @@ -1023,7 +986,7 @@ try [] __device__(T x) -> T { using opmath_t = opmath_type; const opmath_t x_acc = static_cast(x); - return x_acc * compat::tanh(compat::log1p(compat::exp(x_acc))); + return x_acc * ::tanhf(::log1pf(::expf(x_acc))); }, s); From 8774d9891b05865502360bb60b2fde68d8650bd2 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Thu, 9 Jul 2026 13:22:03 +0200 Subject: [PATCH 03/14] nan --- .../bench/transform/applications/P1/pytorch.cu | 13 +++---------- 1 file changed, 3 insertions(+), 10 deletions(-) diff --git a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu index d44cfb58b80..28249e1d5f8 100644 --- a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu +++ b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu @@ -290,13 +290,6 @@ using opmath_type = typename opmath_type_impl::type; // Each wrapper is annotated with the ATen source it replicates. // ============================================================================ -// at::_isnan (ATen/NumericUtils.h:30) — on CUDA this is just ::isnan -template -__device__ __forceinline__ bool _isnan(T val) -{ - return ::isnan(static_cast(val)); -} - // c10::div_floor_floating (c10/util/generic_math.h:34) template __device__ __forceinline__ scalar_t div_floor_floating(scalar_t a, scalar_t b) @@ -638,7 +631,7 @@ try d_b, n, [] __device__(T x, T y) -> T { - if (_isnan(y)) + if (::isnan(static_cast(y))) { return NAN; } @@ -656,7 +649,7 @@ try d_a, n, [] __device__(T x, T y) -> T { - if (_isnan(y)) + if (::isnan(static_cast(y))) { return NAN; } @@ -1109,7 +1102,7 @@ try n, [=] __device__(T v) -> T { using opmath_t = opmath_type; - if (_isnan(static_cast(v))) + if (::isnan(static_cast(v))) { return v; } From 52d39276fe4942fbc20a76c0bbb23ce3c37b3b96 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 10 Jul 2026 11:04:42 +0200 Subject: [PATCH 04/14] hardcode 64 bit offset and add const --- .../transform/applications/P1/pytorch.cu | 115 +++++++++--------- 1 file changed, 55 insertions(+), 60 deletions(-) diff --git a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu index 28249e1d5f8..2602ef18b05 100644 --- a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu +++ b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu @@ -483,8 +483,8 @@ template struct normal_gen { float mean, stddev; - int offset; - __host__ __device__ T operator()(int idx) const + int64_t offset; + __host__ __device__ T operator()(int64_t idx) const { thrust::default_random_engine rng; thrust::random::normal_distribution dist(mean, stddev); @@ -493,14 +493,14 @@ struct normal_gen } }; -template -void fill_normal(thrust::device_vector& v, OffsetT n, int buf_idx) +template +void fill_normal(thrust::device_vector& v, int64_t n, int buf_idx) { v.resize(n); - thrust::transform(thrust::counting_iterator(0), - thrust::counting_iterator(n), + thrust::transform(thrust::counting_iterator(0), + thrust::counting_iterator(n), v.begin(), - normal_gen{0.0f, 1.0f, buf_idx * static_cast(n)}); + normal_gen{0.0f, 1.0f, buf_idx * n}); } // ============================================================================ @@ -508,7 +508,7 @@ void fill_normal(thrust::device_vector& v, OffsetT n, int buf_idx) // ============================================================================ template -void transform(::cuda::std::tuple inputs, Output output, int n, TransformOp op, cudaStream_t stream) +void transform(::cuda::std::tuple inputs, Output output, int64_t n, TransformOp op, cudaStream_t stream) { auto env = cuda::std::execution::env{ ::cuda::stream_ref{stream} @@ -521,7 +521,7 @@ void transform(::cuda::std::tuple inputs, Output output, int n, Trans } template -void transform(Input input, Output output, int n, TransformOp op, cudaStream_t stream) +void transform(Input input, Output output, int64_t n, TransformOp op, cudaStream_t stream) { transform(::cuda::std::make_tuple(input), output, n, op, stream); } @@ -536,11 +536,6 @@ using element_types = nvbench::type_list; using element_types = nvbench::type_list; #endif -// note: pytorch always uses int for offset type -// but be careful not to use too large of a range for the Elements axis, -// otherwise, we will run into overflows. -using pytorch_offset_types = nvbench::type_list; - // ============================================================================ // chained_eltwise_many_in_many_inst // @@ -549,11 +544,11 @@ using pytorch_offset_types = nvbench::type_list; // 11 inputs, 10 binary ops // ============================================================================ -template -static void chained_eltwise_many_in_many_inst(nvbench::state& state, nvbench::type_list) +template +static void chained_eltwise_many_in_many_inst(nvbench::state& state, nvbench::type_list) try { - const auto n = static_cast(state.get_int64("Elements{io}")); + const auto n = state.get_int64("Elements{io}"); constexpr int num_in = 11; thrust::device_vector in[num_in]; @@ -561,7 +556,7 @@ try { fill_normal(in[i], n, i); } - thrust::device_vector tmpA(n), tmpB(n); + thrust::device_vector tmpA(n, thrust::no_init), tmpB(n, thrust::no_init); T* d_in[num_in]; for (int i = 0; i < num_in; i++) @@ -572,15 +567,15 @@ try T* d_b = thrust::raw_pointer_cast(tmpB.data()); state.add_element_count(n); - state.add_global_memory_reads(20L * int64_t{n}); - state.add_global_memory_writes(10L * int64_t{n}); + state.add_global_memory_reads(20L * n); + state.add_global_memory_writes(10L * n); // logaddexp2 captures inv_log_2 — native/cuda/LogAddExpKernel.cu:272 using opmath_t = opmath_type; const auto inv_log_2 = static_cast(1.0 / 0.693147180559945309417232121458176); state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch, [&](const nvbench::launch& launch) { - auto s = launch.get_stream().get_stream(); + const auto s = launch.get_stream().get_stream(); // div_floor: native/cuda/BinaryDivFloorKernel.cu:72, helper c10/util/generic_math.h:34 transform( @@ -727,11 +722,11 @@ catch (const std::bad_alloc&) // 13 inputs, 8 binary ops + 2 ternary ops // ============================================================================ -template -static void chained_eltwise_many_in_few_inst(nvbench::state& state, nvbench::type_list) +template +static void chained_eltwise_many_in_few_inst(nvbench::state& state, nvbench::type_list) try { - const auto n = static_cast(state.get_int64("Elements{io}")); + const auto n = state.get_int64("Elements{io}"); constexpr int num_in = 13; thrust::device_vector in[num_in]; @@ -739,7 +734,7 @@ try { fill_normal(in[i], n, i); } - thrust::device_vector tmpA(n), tmpB(n); + thrust::device_vector tmpA(n, thrust::no_init), tmpB(n, thrust::no_init); T* d_in[num_in]; for (int i = 0; i < num_in; i++) @@ -751,20 +746,20 @@ try // 8 binary (16 reads) + 2 ternary (6 reads) = 22 reads, 10 writes state.add_element_count(n); - state.add_global_memory_reads(22L * int64_t{n}); - state.add_global_memory_writes(10L * int64_t{n}); + state.add_global_memory_reads(22L * n); + state.add_global_memory_writes(10L * n); // Captured scalar parameters, matching how ATen sets them up before gpu_kernel using opmath_t = opmath_type; T beta_val(1.0); // smooth_l1: scalar_t beta_val(beta) T delta_val(1.0); // huber: scalar_t delta_val(delta) // note: opmath_type is same as at::acc_type here - using accscalar_t = opmath_type; // addcmul: at::acc_type - auto alpha = accscalar_t(1); // addcmul: value.to() - auto weight_val = opmath_t(4.0); // lerp scalar: weight.to() + using accscalar_t = opmath_type; // addcmul: at::acc_type + const auto alpha = accscalar_t(1); // addcmul: value.to() + const auto weight_val = opmath_t(4.0); // lerp scalar: weight.to() state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch, [&](const nvbench::launch& launch) { - auto s = launch.get_stream().get_stream(); + const auto s = launch.get_stream().get_stream(); // mse_loss: native/cuda/BinaryMiscOpsKernels.cu:37 transform( @@ -877,35 +872,35 @@ catch (const std::bad_alloc&) // 1 input, 10 unary ops // ============================================================================ -template -static void chained_eltwise_few_in_many_inst(nvbench::state& state, nvbench::type_list) +template +static void chained_eltwise_few_in_many_inst(nvbench::state& state, nvbench::type_list) try { - const auto n = static_cast(state.get_int64("Elements{io}")); + const auto n = state.get_int64("Elements{io}"); - thrust::device_vector input(n); + thrust::device_vector input(n, thrust::no_init); fill_normal(input, n, 0); - thrust::device_vector tmpA(n), tmpB(n); + thrust::device_vector tmpA(n, thrust::no_init), tmpB(n, thrust::no_init); T* d_in = thrust::raw_pointer_cast(input.data()); T* d_a = thrust::raw_pointer_cast(tmpA.data()); T* d_b = thrust::raw_pointer_cast(tmpB.data()); state.add_element_count(n); - state.add_global_memory_reads(10L * int64_t{n}); - state.add_global_memory_writes(10L * int64_t{n}); + state.add_global_memory_reads(10L * n); + state.add_global_memory_writes(10L * n); // Captured scalar parameters - using opmath_t = opmath_type; - const auto exp_val = T(2.5); // pow: exp_scalar.to() - auto beta = opmath_t(1); // softplus: beta_.to() - auto threshold = opmath_t(20); // softplus: threshold_.to() - auto negcoef = opmath_t(1) * opmath_t(1); // elu: alpha * scale - auto poscoef = opmath_t(1); // elu: scale - auto negiptcoef = opmath_t(1); // elu: input_scale + using opmath_t = opmath_type; + const auto exp_val = T(2.5); // pow: exp_scalar.to() + const auto beta = opmath_t(1); // softplus: beta_.to() + const auto threshold = opmath_t(20); // softplus: threshold_.to() + const auto negcoef = opmath_t(1) * opmath_t(1); // elu: alpha * scale + const auto poscoef = opmath_t(1); // elu: scale + const auto negiptcoef = opmath_t(1); // elu: input_scale state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch, [&](const nvbench::launch& launch) { - auto s = launch.get_stream().get_stream(); + const auto s = launch.get_stream().get_stream(); // pow(scalar=2.5): native/cuda/PowKernel.cu:163, helper native/cuda/Pow.cuh:40 transform( @@ -1035,23 +1030,23 @@ catch (const std::bad_alloc&) // 1 input, 10 unary ops // ============================================================================ -template -static void chained_eltwise_few_in_few_inst(nvbench::state& state, nvbench::type_list) +template +static void chained_eltwise_few_in_few_inst(nvbench::state& state, nvbench::type_list) try { - const auto n = static_cast(state.get_int64("Elements{io}")); + const auto n = state.get_int64("Elements{io}"); - thrust::device_vector input(n); + thrust::device_vector input(n, thrust::no_init); fill_normal(input, n, 0); - thrust::device_vector tmpA(n), tmpB(n); + thrust::device_vector tmpA(n, thrust::no_init), tmpB(n, thrust::no_init); T* d_in = thrust::raw_pointer_cast(input.data()); T* d_a = thrust::raw_pointer_cast(tmpA.data()); T* d_b = thrust::raw_pointer_cast(tmpB.data()); state.add_element_count(n); - state.add_global_memory_reads(10L * int64_t{n}); - state.add_global_memory_writes(10L * int64_t{n}); + state.add_global_memory_reads(10L * n); + state.add_global_memory_writes(10L * n); // Captured scalar parameters using opmath_t = opmath_type; @@ -1065,7 +1060,7 @@ try const auto mul_scalar = opmath_t(1.5); // leaky_relu: native/cuda/ActivationLeakyReluKernel.cu:31 - auto negval = opmath_t(0.01); // negval_.to() + const auto negval = opmath_t(0.01); // negval_.to() // hardswish: native/cuda/ActivationHardswishKernel.cu:25 const opmath_t zero(0.0f); @@ -1074,13 +1069,13 @@ try const opmath_t six(6.0f); // hardshrink: native/cuda/ActivationHardshrinkKernel.cu:29 - auto lambd = T(0.5); // value.to() + const auto lambd = T(0.5); // value.to() // gt scalar: native/cuda/CompareKernels.cu:47 const T rhs(0); state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch, [&](const nvbench::launch& launch) { - auto s = launch.get_stream().get_stream(); + const auto s = launch.get_stream().get_stream(); // add(scalar, alpha=1): native/ufunc/add.h:14 transform(d_in, d_a, n, CUDAFunctorOnSelf_add(T(0.5), T(1)), s); @@ -1197,22 +1192,22 @@ catch (const std::bad_alloc&) state.skip("Skipping: out of memory."); } -NVBENCH_BENCH_TYPES(chained_eltwise_many_in_many_inst, NVBENCH_TYPE_AXES(element_types, pytorch_offset_types)) +NVBENCH_BENCH_TYPES(chained_eltwise_many_in_many_inst, NVBENCH_TYPE_AXES(element_types)) .set_name("chained_eltwise_many_in_many_inst") .set_type_axes_names({"T{ct}", "OffsetT{ct}"}) .add_int64_power_of_two_axis("Elements{io}", nvbench::range(16, 28, 4)); -NVBENCH_BENCH_TYPES(chained_eltwise_many_in_few_inst, NVBENCH_TYPE_AXES(element_types, pytorch_offset_types)) +NVBENCH_BENCH_TYPES(chained_eltwise_many_in_few_inst, NVBENCH_TYPE_AXES(element_types)) .set_name("chained_eltwise_many_in_few_inst") .set_type_axes_names({"T{ct}", "OffsetT{ct}"}) .add_int64_power_of_two_axis("Elements{io}", nvbench::range(16, 28, 4)); -NVBENCH_BENCH_TYPES(chained_eltwise_few_in_many_inst, NVBENCH_TYPE_AXES(element_types, pytorch_offset_types)) +NVBENCH_BENCH_TYPES(chained_eltwise_few_in_many_inst, NVBENCH_TYPE_AXES(element_types)) .set_name("chained_eltwise_few_in_many_inst") .set_type_axes_names({"T{ct}", "OffsetT{ct}"}) .add_int64_power_of_two_axis("Elements{io}", nvbench::range(16, 28, 4)); -NVBENCH_BENCH_TYPES(chained_eltwise_few_in_few_inst, NVBENCH_TYPE_AXES(element_types, pytorch_offset_types)) +NVBENCH_BENCH_TYPES(chained_eltwise_few_in_few_inst, NVBENCH_TYPE_AXES(element_types)) .set_name("chained_eltwise_few_in_few_inst") .set_type_axes_names({"T{ct}", "OffsetT{ct}"}) .add_int64_power_of_two_axis("Elements{io}", nvbench::range(16, 28, 4)); From ddc24aa8b5cd7e7a6f962f39e37c43f3c55b0b90 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 10 Jul 2026 11:07:06 +0200 Subject: [PATCH 05/14] namesopace --- .../transform/applications/P1/pytorch.cu | 46 +++++++++---------- 1 file changed, 23 insertions(+), 23 deletions(-) diff --git a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu index 2602ef18b05..3d05293bdec 100644 --- a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu +++ b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu @@ -508,10 +508,10 @@ void fill_normal(thrust::device_vector& v, int64_t n, int buf_idx) // ============================================================================ template -void transform(::cuda::std::tuple inputs, Output output, int64_t n, TransformOp op, cudaStream_t stream) +void transform(cuda::std::tuple inputs, Output output, int64_t n, TransformOp op, cudaStream_t stream) { auto env = cuda::std::execution::env{ - ::cuda::stream_ref{stream} + cuda::stream_ref{stream} #if !TUNE_BASE , cuda::execution::tune(policy_selector{}) @@ -523,7 +523,7 @@ void transform(::cuda::std::tuple inputs, Output output, int64_t n, T template void transform(Input input, Output output, int64_t n, TransformOp op, cudaStream_t stream) { - transform(::cuda::std::make_tuple(input), output, n, op, stream); + transform(cuda::std::make_tuple(input), output, n, op, stream); } // ============================================================================ @@ -579,7 +579,7 @@ try // div_floor: native/cuda/BinaryDivFloorKernel.cu:72, helper c10/util/generic_math.h:34 transform( - ::cuda::std::make_tuple(d_in[0], d_in[1]), + cuda::std::make_tuple(d_in[0], d_in[1]), d_a, n, [] __device__(T a, T b) -> T { @@ -589,7 +589,7 @@ try // div_trunc: native/cuda/BinaryDivTruncKernel.cu:42 transform( - ::cuda::std::make_tuple(d_a, d_in[2]), + cuda::std::make_tuple(d_a, d_in[2]), d_b, n, [] __device__(T a, T b) -> T { @@ -598,11 +598,11 @@ try s); // div: native/cuda/BinaryDivTrueKernel.cu:54, DivFunctor in native/cuda/BinaryInternal.h:20 - transform(::cuda::std::make_tuple(d_b, d_in[3]), d_a, n, DivFunctor(), s); + transform(cuda::std::make_tuple(d_b, d_in[3]), d_a, n, DivFunctor(), s); // atan2: native/cuda/BinaryGeometricKernels.cu:18 transform( - ::cuda::std::make_tuple(d_a, d_in[4]), + cuda::std::make_tuple(d_a, d_in[4]), d_b, n, [] __device__(T a, T b) -> T { @@ -612,7 +612,7 @@ try // hypot: native/cuda/BinaryGeometricKernels.cu:29 transform( - ::cuda::std::make_tuple(d_b, d_in[5]), + cuda::std::make_tuple(d_b, d_in[5]), d_a, n, [] __device__(T a, T b) -> T { @@ -622,7 +622,7 @@ try // xlogy: native/cuda/BinaryMiscOpsKernels.cu:46 transform( - ::cuda::std::make_tuple(d_a, d_in[6]), + cuda::std::make_tuple(d_a, d_in[6]), d_b, n, [] __device__(T x, T y) -> T { @@ -640,7 +640,7 @@ try // xlog1py: native/cuda/BinaryMiscOpsKernels.cu:60 transform( - ::cuda::std::make_tuple(d_b, d_in[7]), + cuda::std::make_tuple(d_b, d_in[7]), d_a, n, [] __device__(T x, T y) -> T { @@ -658,7 +658,7 @@ try // logaddexp: native/cuda/LogAddExpKernel.cu:253 transform( - ::cuda::std::make_tuple(d_a, d_in[8]), + cuda::std::make_tuple(d_a, d_in[8]), d_b, n, [] __device__(T a_, T b_) -> T { @@ -679,7 +679,7 @@ try // logaddexp2: native/cuda/LogAddExpKernel.cu:272 transform( - ::cuda::std::make_tuple(d_b, d_in[9]), + cuda::std::make_tuple(d_b, d_in[9]), d_a, n, [inv_log_2] __device__(T a_, T b_) -> T { @@ -700,7 +700,7 @@ try // pow (tensor,tensor): native/cuda/PowKernel.cu:136, helper native/cuda/Pow.cuh:40 transform( - ::cuda::std::make_tuple(d_a, d_in[10]), + cuda::std::make_tuple(d_a, d_in[10]), d_b, n, [] __device__(T base, T exp) -> T { @@ -763,7 +763,7 @@ try // mse_loss: native/cuda/BinaryMiscOpsKernels.cu:37 transform( - ::cuda::std::make_tuple(d_in[0], d_in[1]), + cuda::std::make_tuple(d_in[0], d_in[1]), d_a, n, [] __device__(T a, T b) -> T { @@ -774,7 +774,7 @@ try // smooth_l1_loss(beta=1.0): native/cuda/BinaryMiscOpsKernels.cu:19 transform( - ::cuda::std::make_tuple(d_a, d_in[2]), + cuda::std::make_tuple(d_a, d_in[2]), d_b, n, [beta_val] __device__(T a, T b) -> T { @@ -785,7 +785,7 @@ try // huber_loss(delta=1.0): native/cuda/BinaryMiscOpsKernels.cu:29 transform( - ::cuda::std::make_tuple(d_b, d_in[3]), + cuda::std::make_tuple(d_b, d_in[3]), d_a, n, [delta_val] __device__(T a, T b) -> T { @@ -796,7 +796,7 @@ try // clamp(min=tensor) -> maximum: native/cuda/MaxMinElementwiseKernel.cu:28 transform( - ::cuda::std::make_tuple(d_a, d_in[4]), + cuda::std::make_tuple(d_a, d_in[4]), d_b, n, [] __device__(T a, T b) -> T { @@ -817,14 +817,14 @@ try // mul: native/cuda/BinaryMulKernel.cu:39, MulFunctor in native/cuda/BinaryInternal.h:27 using mul_opmath_t = opmath_type; - transform(::cuda::std::make_tuple(d_b, d_in[5]), d_a, n, MulFunctor(), s); + transform(cuda::std::make_tuple(d_b, d_in[5]), d_a, n, MulFunctor(), s); // add(alpha=1): native/ufunc/add.h:14, torchgen/dest/ufunc.py - transform(::cuda::std::make_tuple(d_a, d_in[6]), d_b, n, CUDAFunctor_add(1.0), s); + transform(cuda::std::make_tuple(d_a, d_in[6]), d_b, n, CUDAFunctor_add(1.0), s); // addcmul(value=1): native/cuda/PointwiseOpsKernel.cu:87, native/cuda/DeviceAddCmulCdiv.cuh:9 transform( - ::cuda::std::make_tuple(d_b, d_in[7], d_in[8]), + cuda::std::make_tuple(d_b, d_in[7], d_in[8]), d_a, n, [alpha] __device__(T a, T b, T c) -> T { @@ -834,7 +834,7 @@ try // lerp(weight=4.0): native/cuda/Lerp.cu:130, native/Lerp.h:21 transform( - ::cuda::std::make_tuple(d_a, d_in[9]), + cuda::std::make_tuple(d_a, d_in[9]), d_b, n, [=] __device__(T self_val, T end_val) { @@ -844,7 +844,7 @@ try // lerp(weight=tensor): native/cuda/Lerp.cu:76, native/Lerp.h:21 transform( - ::cuda::std::make_tuple(d_b, d_in[10], d_in[11]), + cuda::std::make_tuple(d_b, d_in[10], d_in[11]), d_a, n, [] __device__(T self_val, T end_val, T weight_val) -> T { @@ -856,7 +856,7 @@ try // it must hold at least enough memory per element for bool (1 byte) // greater: native/cuda/CompareKernels.cu:69 CompareFunctor comp_f(OpType::GT); - transform(::cuda::std::make_tuple(d_a, d_in[12]), d_b, n, comp_f, s); + transform(cuda::std::make_tuple(d_a, d_in[12]), d_b, n, comp_f, s); }); } catch (const std::bad_alloc&) From ff737ec8762c3fefd7210a31fdfddb3de3dca47c Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 10 Jul 2026 11:09:15 +0200 Subject: [PATCH 06/14] Move bf16 to header --- .../transform/applications/P1/bfloat16.h | 229 +++++++++++++++++ .../transform/applications/P1/pytorch.cu | 232 +----------------- 2 files changed, 230 insertions(+), 231 deletions(-) create mode 100644 cub/benchmarks/bench/transform/applications/P1/bfloat16.h diff --git a/cub/benchmarks/bench/transform/applications/P1/bfloat16.h b/cub/benchmarks/bench/transform/applications/P1/bfloat16.h new file mode 100644 index 00000000000..75b1883197b --- /dev/null +++ b/cub/benchmarks/bench/transform/applications/P1/bfloat16.h @@ -0,0 +1,229 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#pragma once + +#include +#include + +#include + +// ============================================================================ +// BFloat16 type — replicated from c10::BFloat16 +// (torch/headeronly/util/BFloat16.h) +// ============================================================================ + +namespace bf16_detail +{ +inline __host__ __device__ float f32_from_bits(uint16_t src) +{ + float res = 0; + uint32_t tmp = src; + tmp <<= 16; + std::memcpy(&res, &tmp, sizeof(tmp)); + return res; +} + +inline __host__ __device__ uint16_t round_to_nearest_even(float src) +{ + if (std::isnan(src)) + { + return UINT16_C(0x7FC0); + } + else + { + uint32_t U32; + std::memcpy(&U32, &src, sizeof(U32)); + uint32_t rounding_bias = ((U32 >> 16) & 1) + UINT32_C(0x7FFF); + return static_cast((U32 + rounding_bias) >> 16); + } +} +} // namespace bf16_detail + +struct alignas(2) BFloat16 +{ + uint16_t x; + + BFloat16() = default; + + struct from_bits_t + {}; + static constexpr __host__ __device__ from_bits_t from_bits() + { + return from_bits_t(); + } + + constexpr __host__ __device__ BFloat16(unsigned short bits, from_bits_t) + : x(bits) + {} + + /* implicit */ inline __host__ __device__ BFloat16(float value); + inline __host__ __device__ operator float() const; + + inline __host__ __device__ BFloat16(const __nv_bfloat16& value); + explicit inline __host__ __device__ operator __nv_bfloat16() const; +}; + +inline __host__ __device__ BFloat16::BFloat16(float value) +{ + NV_IF_TARGET(NV_IS_DEVICE, + (x = __bfloat16_as_ushort(__float2bfloat16(value));), + (x = bf16_detail::round_to_nearest_even(value);)); +} + +inline __host__ __device__ BFloat16::operator float() const +{ + return __bfloat162float(*reinterpret_cast(&x)); +} + +inline __host__ __device__ BFloat16::BFloat16(const __nv_bfloat16& value) +{ + x = *reinterpret_cast(&value); +} +inline __host__ __device__ BFloat16::operator __nv_bfloat16() const +{ + return *reinterpret_cast(&x); +} + +// Arithmetic — BFloat16 x BFloat16 → BFloat16 +inline __host__ __device__ BFloat16 operator+(const BFloat16& a, const BFloat16& b) +{ + return static_cast(a) + static_cast(b); +} +inline __host__ __device__ BFloat16 operator-(const BFloat16& a, const BFloat16& b) +{ + return static_cast(a) - static_cast(b); +} +inline __host__ __device__ BFloat16 operator*(const BFloat16& a, const BFloat16& b) +{ + return static_cast(a) * static_cast(b); +} +inline __host__ __device__ BFloat16 operator/(const BFloat16& a, const BFloat16& b) +{ + return static_cast(a) / static_cast(b); +} +inline __host__ __device__ BFloat16 operator-(const BFloat16& a) +{ + return -static_cast(a); +} + +// Compound assignment — BFloat16 +inline __host__ __device__ BFloat16& operator+=(BFloat16& a, const BFloat16& b) +{ + a = a + b; + return a; +} +inline __host__ __device__ BFloat16& operator-=(BFloat16& a, const BFloat16& b) +{ + a = a - b; + return a; +} +inline __host__ __device__ BFloat16& operator*=(BFloat16& a, const BFloat16& b) +{ + a = a * b; + return a; +} +inline __host__ __device__ BFloat16& operator/=(BFloat16& a, const BFloat16& b) +{ + a = a / b; + return a; +} + +// Arithmetic — BFloat16 x float → float +inline __host__ __device__ float operator+(BFloat16 a, float b) +{ + return static_cast(a) + b; +} +inline __host__ __device__ float operator-(BFloat16 a, float b) +{ + return static_cast(a) - b; +} +inline __host__ __device__ float operator*(BFloat16 a, float b) +{ + return static_cast(a) * b; +} +inline __host__ __device__ float operator/(BFloat16 a, float b) +{ + return static_cast(a) / b; +} +inline __host__ __device__ float operator+(float a, BFloat16 b) +{ + return a + static_cast(b); +} +inline __host__ __device__ float operator-(float a, BFloat16 b) +{ + return a - static_cast(b); +} +inline __host__ __device__ float operator*(float a, BFloat16 b) +{ + return a * static_cast(b); +} +inline __host__ __device__ float operator/(float a, BFloat16 b) +{ + return a / static_cast(b); +} + +// Compound assignment — float x BFloat16 → float +inline __host__ __device__ float& operator+=(float& a, const BFloat16& b) +{ + return a += static_cast(b); +} +inline __host__ __device__ float& operator-=(float& a, const BFloat16& b) +{ + return a -= static_cast(b); +} +inline __host__ __device__ float& operator*=(float& a, const BFloat16& b) +{ + return a *= static_cast(b); +} +inline __host__ __device__ float& operator/=(float& a, const BFloat16& b) +{ + return a /= static_cast(b); +} + +// Arithmetic — BFloat16 x int → BFloat16 +inline __host__ __device__ BFloat16 operator+(BFloat16 a, int b) +{ + return a + static_cast(static_cast(b)); +} +inline __host__ __device__ BFloat16 operator-(BFloat16 a, int b) +{ + return a - static_cast(static_cast(b)); +} +inline __host__ __device__ BFloat16 operator*(BFloat16 a, int b) +{ + return a * static_cast(static_cast(b)); +} +inline __host__ __device__ BFloat16 operator/(BFloat16 a, int b) +{ + return a / static_cast(static_cast(b)); +} +inline __host__ __device__ BFloat16 operator+(int a, BFloat16 b) +{ + return static_cast(static_cast(a)) + b; +} +inline __host__ __device__ BFloat16 operator-(int a, BFloat16 b) +{ + return static_cast(static_cast(a)) - b; +} +inline __host__ __device__ BFloat16 operator*(int a, BFloat16 b) +{ + return static_cast(static_cast(a)) * b; +} +inline __host__ __device__ BFloat16 operator/(int a, BFloat16 b) +{ + return static_cast(static_cast(a)) / b; +} + +// Comparison — for std::min/std::max +inline __host__ __device__ bool operator>(BFloat16& lhs, BFloat16& rhs) +{ + return float(lhs) > float(rhs); +} +inline __host__ __device__ bool operator<(BFloat16& lhs, BFloat16& rhs) +{ + return float(lhs) < float(rhs); +} + +// NVBench type registration +NVBENCH_DECLARE_TYPE_STRINGS(BFloat16, "bf16", "BFloat16"); diff --git a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu index 3d05293bdec..4b128f17398 100644 --- a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu +++ b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu @@ -31,241 +31,11 @@ #include #include -#include #include #include #include "../../common.h" - -// ============================================================================ -// BFloat16 type — replicated from c10::BFloat16 -// (torch/headeronly/util/BFloat16.h) -// ============================================================================ - -namespace bf16_detail -{ -inline __host__ __device__ float f32_from_bits(uint16_t src) -{ - float res = 0; - uint32_t tmp = src; - tmp <<= 16; - std::memcpy(&res, &tmp, sizeof(tmp)); - return res; -} - -inline __host__ __device__ uint16_t round_to_nearest_even(float src) -{ - if (std::isnan(src)) - { - return UINT16_C(0x7FC0); - } - else - { - uint32_t U32; - std::memcpy(&U32, &src, sizeof(U32)); - uint32_t rounding_bias = ((U32 >> 16) & 1) + UINT32_C(0x7FFF); - return static_cast((U32 + rounding_bias) >> 16); - } -} -} // namespace bf16_detail - -struct alignas(2) BFloat16 -{ - uint16_t x; - - BFloat16() = default; - - struct from_bits_t - {}; - static constexpr __host__ __device__ from_bits_t from_bits() - { - return from_bits_t(); - } - - constexpr __host__ __device__ BFloat16(unsigned short bits, from_bits_t) - : x(bits) - {} - - /* implicit */ inline __host__ __device__ BFloat16(float value); - inline __host__ __device__ operator float() const; - -#if defined(__CUDACC__) - inline __host__ __device__ BFloat16(const __nv_bfloat16& value); - explicit inline __host__ __device__ operator __nv_bfloat16() const; -#endif -}; - -inline __host__ __device__ BFloat16::BFloat16(float value) - : -#if defined(__CUDACC__) && defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800 - x(__bfloat16_as_ushort(__float2bfloat16(value))) -#else - x(bf16_detail::round_to_nearest_even(value)) -#endif -{} - -inline __host__ __device__ BFloat16::operator float() const -{ -#if defined(__CUDACC__) - return __bfloat162float(*reinterpret_cast(&x)); -#else - return bf16_detail::f32_from_bits(x); -#endif -} - -#if defined(__CUDACC__) -inline __host__ __device__ BFloat16::BFloat16(const __nv_bfloat16& value) -{ - x = *reinterpret_cast(&value); -} -inline __host__ __device__ BFloat16::operator __nv_bfloat16() const -{ - return *reinterpret_cast(&x); -} -#endif - -// Arithmetic — BFloat16 x BFloat16 → BFloat16 -inline __host__ __device__ BFloat16 operator+(const BFloat16& a, const BFloat16& b) -{ - return static_cast(a) + static_cast(b); -} -inline __host__ __device__ BFloat16 operator-(const BFloat16& a, const BFloat16& b) -{ - return static_cast(a) - static_cast(b); -} -inline __host__ __device__ BFloat16 operator*(const BFloat16& a, const BFloat16& b) -{ - return static_cast(a) * static_cast(b); -} -inline __host__ __device__ BFloat16 operator/(const BFloat16& a, const BFloat16& b) -{ - return static_cast(a) / static_cast(b); -} -inline __host__ __device__ BFloat16 operator-(const BFloat16& a) -{ - return -static_cast(a); -} - -// Compound assignment — BFloat16 -inline __host__ __device__ BFloat16& operator+=(BFloat16& a, const BFloat16& b) -{ - a = a + b; - return a; -} -inline __host__ __device__ BFloat16& operator-=(BFloat16& a, const BFloat16& b) -{ - a = a - b; - return a; -} -inline __host__ __device__ BFloat16& operator*=(BFloat16& a, const BFloat16& b) -{ - a = a * b; - return a; -} -inline __host__ __device__ BFloat16& operator/=(BFloat16& a, const BFloat16& b) -{ - a = a / b; - return a; -} - -// Arithmetic — BFloat16 x float → float -inline __host__ __device__ float operator+(BFloat16 a, float b) -{ - return static_cast(a) + b; -} -inline __host__ __device__ float operator-(BFloat16 a, float b) -{ - return static_cast(a) - b; -} -inline __host__ __device__ float operator*(BFloat16 a, float b) -{ - return static_cast(a) * b; -} -inline __host__ __device__ float operator/(BFloat16 a, float b) -{ - return static_cast(a) / b; -} -inline __host__ __device__ float operator+(float a, BFloat16 b) -{ - return a + static_cast(b); -} -inline __host__ __device__ float operator-(float a, BFloat16 b) -{ - return a - static_cast(b); -} -inline __host__ __device__ float operator*(float a, BFloat16 b) -{ - return a * static_cast(b); -} -inline __host__ __device__ float operator/(float a, BFloat16 b) -{ - return a / static_cast(b); -} - -// Compound assignment — float x BFloat16 → float -inline __host__ __device__ float& operator+=(float& a, const BFloat16& b) -{ - return a += static_cast(b); -} -inline __host__ __device__ float& operator-=(float& a, const BFloat16& b) -{ - return a -= static_cast(b); -} -inline __host__ __device__ float& operator*=(float& a, const BFloat16& b) -{ - return a *= static_cast(b); -} -inline __host__ __device__ float& operator/=(float& a, const BFloat16& b) -{ - return a /= static_cast(b); -} - -// Arithmetic — BFloat16 x int → BFloat16 -inline __host__ __device__ BFloat16 operator+(BFloat16 a, int b) -{ - return a + static_cast(static_cast(b)); -} -inline __host__ __device__ BFloat16 operator-(BFloat16 a, int b) -{ - return a - static_cast(static_cast(b)); -} -inline __host__ __device__ BFloat16 operator*(BFloat16 a, int b) -{ - return a * static_cast(static_cast(b)); -} -inline __host__ __device__ BFloat16 operator/(BFloat16 a, int b) -{ - return a / static_cast(static_cast(b)); -} -inline __host__ __device__ BFloat16 operator+(int a, BFloat16 b) -{ - return static_cast(static_cast(a)) + b; -} -inline __host__ __device__ BFloat16 operator-(int a, BFloat16 b) -{ - return static_cast(static_cast(a)) - b; -} -inline __host__ __device__ BFloat16 operator*(int a, BFloat16 b) -{ - return static_cast(static_cast(a)) * b; -} -inline __host__ __device__ BFloat16 operator/(int a, BFloat16 b) -{ - return static_cast(static_cast(a)) / b; -} - -// Comparison — for std::min/std::max -inline __host__ __device__ bool operator>(BFloat16& lhs, BFloat16& rhs) -{ - return float(lhs) > float(rhs); -} -inline __host__ __device__ bool operator<(BFloat16& lhs, BFloat16& rhs) -{ - return float(lhs) < float(rhs); -} - -// NVBench type registration -NVBENCH_DECLARE_TYPE_STRINGS(BFloat16, "bf16", "BFloat16"); +#include "bfloat16.h" // ============================================================================ // at::opmath_type — the compute type for intermediate math From 84b09ac4c643936d4437229d1f6cf898385f3306 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 10 Jul 2026 11:13:30 +0200 Subject: [PATCH 07/14] Remove pow --- .../bench/transform/applications/P1/pytorch.cu | 11 ++--------- 1 file changed, 2 insertions(+), 9 deletions(-) diff --git a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu index 4b128f17398..4805f0a5999 100644 --- a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu +++ b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu @@ -92,13 +92,6 @@ __device__ __forceinline__ scalar_t div_floor_floating(scalar_t a, scalar_t b) return floordiv; } -// pow_ (native/cuda/Pow.cuh:40) -template -__device__ __forceinline__ Base_type pow_(Base_type base, Exp_type exp) -{ - return ::pow(base, exp); -} - // is_lerp_weight_small + lerp (native/Lerp.h:11,21) template __device__ __forceinline__ bool is_lerp_weight_small(scalar_t weight) @@ -474,7 +467,7 @@ try d_b, n, [] __device__(T base, T exp) -> T { - return pow_(base, exp); + return cuda::std::pow(base, exp); }, s); }); @@ -678,7 +671,7 @@ try d_a, n, [=] __device__(T base) -> T { - return pow_(base, exp_val); + return cuda::std::pow(base, exp_val); }, s); From 63f97a28728b898ab79ca2e5efb8525b8ece4175 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 10 Jul 2026 11:23:36 +0200 Subject: [PATCH 08/14] Document benchmark --- .../transform/applications/P1/pytorch.cu | 41 +++++++++++-------- 1 file changed, 23 insertions(+), 18 deletions(-) diff --git a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu index 4805f0a5999..0ef5c98f0f1 100644 --- a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu +++ b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu @@ -1,8 +1,13 @@ // SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// NVBench benchmarks for chained elementwise operations put together by Matthias Jouanneaux (DevTech). -// See also: https://github.com/NVIDIA-dev/cccl_private/issues/639 +// ============================================================================ +// NVBench benchmarks for chained elementwise operations was put together by Matthias Jouanneaux (DevTech). It mimics +// how pytorch uses element-wise kernels (i.e. cub::DeviceTransform). The main difference to ordinary CCCL benchmarks is +// the chaining of several operations in the benchmark's critical section. Furthermore, there are 4 types of work loads +// covering the combinations of few vs. many input buffers and few vs. many instructions in the kernel. +// For more discussion see: https://github.com/NVIDIA-dev/cccl_private/issues/639 +// ============================================================================ // %RANGE% TUNE_BIF_BIAS bif -16:16:4 // %RANGE% TUNE_ALGORITHM alg 0:4:1 @@ -300,7 +305,7 @@ using element_types = nvbench::type_list; #endif // ============================================================================ -// chained_eltwise_many_in_many_inst +// many_inputs_many_instructions // // div_floor -> div_trunc -> div -> atan2 -> hypot -> // xlogy -> xlog1py -> logaddexp -> logaddexp2 -> pow @@ -308,7 +313,7 @@ using element_types = nvbench::type_list; // ============================================================================ template -static void chained_eltwise_many_in_many_inst(nvbench::state& state, nvbench::type_list) +static void many_inputs_many_instructions(nvbench::state& state, nvbench::type_list) try { const auto n = state.get_int64("Elements{io}"); @@ -478,7 +483,7 @@ catch (const std::bad_alloc&) } // ============================================================================ -// chained_eltwise_many_in_few_inst +// many_inputs_few_instructions // // mse_loss -> smooth_l1_loss -> huber_loss -> clamp_min -> // mul -> add -> addcmul -> lerp(scalar) -> lerp(tensor) -> greater @@ -486,7 +491,7 @@ catch (const std::bad_alloc&) // ============================================================================ template -static void chained_eltwise_many_in_few_inst(nvbench::state& state, nvbench::type_list) +static void many_inputs_few_instructions(nvbench::state& state, nvbench::type_list) try { const auto n = state.get_int64("Elements{io}"); @@ -628,7 +633,7 @@ catch (const std::bad_alloc&) } // ============================================================================ -// chained_eltwise_few_in_many_inst +// few_inputs_many_instructions // // pow(2.5) -> tanh -> sin -> cos -> softplus -> // silu -> mish -> elu -> gelu -> logsigmoid @@ -636,7 +641,7 @@ catch (const std::bad_alloc&) // ============================================================================ template -static void chained_eltwise_few_in_many_inst(nvbench::state& state, nvbench::type_list) +static void few_inputs_many_instructions(nvbench::state& state, nvbench::type_list) try { const auto n = state.get_int64("Elements{io}"); @@ -786,7 +791,7 @@ catch (const std::bad_alloc&) } // ============================================================================ -// chained_eltwise_few_in_few_inst +// few_inputs_few_instructions // // add(0.5) -> neg -> clamp(-2,1) -> abs -> mul(1.5) -> // leaky_relu -> hardswish -> hardshrink -> hardsigmoid -> gt(0) @@ -794,7 +799,7 @@ catch (const std::bad_alloc&) // ============================================================================ template -static void chained_eltwise_few_in_few_inst(nvbench::state& state, nvbench::type_list) +static void few_inputs_few_instructions(nvbench::state& state, nvbench::type_list) try { const auto n = state.get_int64("Elements{io}"); @@ -955,22 +960,22 @@ catch (const std::bad_alloc&) state.skip("Skipping: out of memory."); } -NVBENCH_BENCH_TYPES(chained_eltwise_many_in_many_inst, NVBENCH_TYPE_AXES(element_types)) - .set_name("chained_eltwise_many_in_many_inst") +NVBENCH_BENCH_TYPES(many_inputs_many_instructions, NVBENCH_TYPE_AXES(element_types)) + .set_name("many_inputs_many_instructions") .set_type_axes_names({"T{ct}", "OffsetT{ct}"}) .add_int64_power_of_two_axis("Elements{io}", nvbench::range(16, 28, 4)); -NVBENCH_BENCH_TYPES(chained_eltwise_many_in_few_inst, NVBENCH_TYPE_AXES(element_types)) - .set_name("chained_eltwise_many_in_few_inst") +NVBENCH_BENCH_TYPES(many_inputs_few_instructions, NVBENCH_TYPE_AXES(element_types)) + .set_name("many_inputs_few_instructions") .set_type_axes_names({"T{ct}", "OffsetT{ct}"}) .add_int64_power_of_two_axis("Elements{io}", nvbench::range(16, 28, 4)); -NVBENCH_BENCH_TYPES(chained_eltwise_few_in_many_inst, NVBENCH_TYPE_AXES(element_types)) - .set_name("chained_eltwise_few_in_many_inst") +NVBENCH_BENCH_TYPES(few_inputs_many_instructions, NVBENCH_TYPE_AXES(element_types)) + .set_name("few_inputs_many_instructions") .set_type_axes_names({"T{ct}", "OffsetT{ct}"}) .add_int64_power_of_two_axis("Elements{io}", nvbench::range(16, 28, 4)); -NVBENCH_BENCH_TYPES(chained_eltwise_few_in_few_inst, NVBENCH_TYPE_AXES(element_types)) - .set_name("chained_eltwise_few_in_few_inst") +NVBENCH_BENCH_TYPES(few_inputs_few_instructions, NVBENCH_TYPE_AXES(element_types)) + .set_name("few_inputs_few_instructions") .set_type_axes_names({"T{ct}", "OffsetT{ct}"}) .add_int64_power_of_two_axis("Elements{io}", nvbench::range(16, 28, 4)); From 3cca58a900713cbc10dddca4fcd2a5729b4f63d9 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 10 Jul 2026 11:30:42 +0200 Subject: [PATCH 09/14] Document benchmark --- .../bench/transform/applications/P1/pytorch.cu | 9 +++++---- 1 file changed, 5 insertions(+), 4 deletions(-) diff --git a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu index 0ef5c98f0f1..140dc3a874f 100644 --- a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu +++ b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu @@ -3,10 +3,11 @@ // ============================================================================ // NVBench benchmarks for chained elementwise operations was put together by Matthias Jouanneaux (DevTech). It mimics -// how pytorch uses element-wise kernels (i.e. cub::DeviceTransform). The main difference to ordinary CCCL benchmarks is -// the chaining of several operations in the benchmark's critical section. Furthermore, there are 4 types of work loads -// covering the combinations of few vs. many input buffers and few vs. many instructions in the kernel. -// For more discussion see: https://github.com/NVIDIA-dev/cccl_private/issues/639 +// how pytorch uses element-wise kernels (i.e. cub::DeviceTransform) and also tries to preserve pytorch's operators and +// utility types. The main difference to ordinary CCCL benchmarks is the chaining of several operations in the +// benchmark's critical section. Furthermore, there are 4 types of work loads covering the combinations of few vs. many +// input buffers and few vs. many instructions in the kernel. For more discussion see: +// https://github.com/NVIDIA-dev/cccl_private/issues/639 // ============================================================================ // %RANGE% TUNE_BIF_BIAS bif -16:16:4 From 3275ba50fb9ce4217411e378583ab9cb36daa5ec Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 10 Jul 2026 11:45:21 +0200 Subject: [PATCH 10/14] Replace more stuff by cuda --- .../transform/applications/P1/pytorch.cu | 27 ++++++++++--------- 1 file changed, 15 insertions(+), 12 deletions(-) diff --git a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu index 140dc3a874f..6e9042050e3 100644 --- a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu +++ b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu @@ -32,13 +32,14 @@ #endif // !TUNE_BASE && TUNE_ALGORITHM != 1 && (TUNE_VEC_SIZE_POW2 != 1) #include -#include -#include -#include -#include -#include -#include +#include +#include +#include +#include +#include +#include +#include #include "../../common.h" #include "bfloat16.h" @@ -255,9 +256,9 @@ struct normal_gen int64_t offset; __host__ __device__ T operator()(int64_t idx) const { - thrust::default_random_engine rng; - thrust::random::normal_distribution dist(mean, stddev); + cuda::pcg64 rng(42); rng.discard(offset + idx); + cuda::std::normal_distribution dist(mean, stddev); return T(dist(rng)); } }; @@ -266,10 +267,12 @@ template void fill_normal(thrust::device_vector& v, int64_t n, int buf_idx) { v.resize(n); - thrust::transform(thrust::counting_iterator(0), - thrust::counting_iterator(n), - v.begin(), - normal_gen{0.0f, 1.0f, buf_idx * n}); + cuda::std::transform( + cuda::execution::gpu, + cuda::counting_iterator(0), + cuda::counting_iterator(n), + v.begin(), + normal_gen{0.0f, 1.0f, buf_idx * n}); } // ============================================================================ From cdbb19f4fe5f966068a19da12a76006d78819e62 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 10 Jul 2026 11:45:44 +0200 Subject: [PATCH 11/14] Drop offset description --- cub/benchmarks/bench/transform/applications/P1/pytorch.cu | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu index 6e9042050e3..c97bc0eb945 100644 --- a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu +++ b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu @@ -966,20 +966,20 @@ catch (const std::bad_alloc&) NVBENCH_BENCH_TYPES(many_inputs_many_instructions, NVBENCH_TYPE_AXES(element_types)) .set_name("many_inputs_many_instructions") - .set_type_axes_names({"T{ct}", "OffsetT{ct}"}) + .set_type_axes_names({"T{ct}"}) .add_int64_power_of_two_axis("Elements{io}", nvbench::range(16, 28, 4)); NVBENCH_BENCH_TYPES(many_inputs_few_instructions, NVBENCH_TYPE_AXES(element_types)) .set_name("many_inputs_few_instructions") - .set_type_axes_names({"T{ct}", "OffsetT{ct}"}) + .set_type_axes_names({"T{ct}"}) .add_int64_power_of_two_axis("Elements{io}", nvbench::range(16, 28, 4)); NVBENCH_BENCH_TYPES(few_inputs_many_instructions, NVBENCH_TYPE_AXES(element_types)) .set_name("few_inputs_many_instructions") - .set_type_axes_names({"T{ct}", "OffsetT{ct}"}) + .set_type_axes_names({"T{ct}"}) .add_int64_power_of_two_axis("Elements{io}", nvbench::range(16, 28, 4)); NVBENCH_BENCH_TYPES(few_inputs_few_instructions, NVBENCH_TYPE_AXES(element_types)) .set_name("few_inputs_few_instructions") - .set_type_axes_names({"T{ct}", "OffsetT{ct}"}) + .set_type_axes_names({"T{ct}"}) .add_int64_power_of_two_axis("Elements{io}", nvbench::range(16, 28, 4)); From 0e5b684ce183ca8216a8ac42e317b162176a7bc8 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 10 Jul 2026 15:10:41 +0200 Subject: [PATCH 12/14] Replace __bfloat16_as_ushort by reinterpret_cast --- .../bench/transform/applications/P1/bfloat16.h | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/cub/benchmarks/bench/transform/applications/P1/bfloat16.h b/cub/benchmarks/bench/transform/applications/P1/bfloat16.h index 75b1883197b..275fcfe5dcc 100644 --- a/cub/benchmarks/bench/transform/applications/P1/bfloat16.h +++ b/cub/benchmarks/bench/transform/applications/P1/bfloat16.h @@ -3,7 +3,10 @@ #pragma once +#include + #include +#include #include #include @@ -67,8 +70,11 @@ struct alignas(2) BFloat16 inline __host__ __device__ BFloat16::BFloat16(float value) { NV_IF_TARGET(NV_IS_DEVICE, - (x = __bfloat16_as_ushort(__float2bfloat16(value));), - (x = bf16_detail::round_to_nearest_even(value);)); + ({ + __nv_bfloat16 tmp = __float2bfloat16(value); + x = *reinterpret_cast(&tmp); + }), + ({ x = bf16_detail::round_to_nearest_even(value); })); } inline __host__ __device__ BFloat16::operator float() const From 607208876be6c7cfa85cb87d81fb119af07bf3bb Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 10 Jul 2026 15:49:50 +0200 Subject: [PATCH 13/14] Drop --expt-relaxed-constexpr --- cub/benchmarks/CMakeLists.txt | 3 +-- .../bench/transform/applications/P1/pytorch.cu | 10 ++++++---- 2 files changed, 7 insertions(+), 6 deletions(-) diff --git a/cub/benchmarks/CMakeLists.txt b/cub/benchmarks/CMakeLists.txt index 321c81e26e2..72fcdd36b50 100644 --- a/cub/benchmarks/CMakeLists.txt +++ b/cub/benchmarks/CMakeLists.txt @@ -124,8 +124,7 @@ function(add_bench_dir bench_dir) target_compile_definitions(${base_bench_target} PRIVATE TUNE_BASE=1) target_compile_options( ${base_bench_target} - PRIVATE - "$<$:--extended-lambda;--expt-relaxed-constexpr>" + PRIVATE "$<$:--extended-lambda>" ) if (CUB_ENABLE_TUNING) diff --git a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu index c97bc0eb945..4c379f0587a 100644 --- a/cub/benchmarks/bench/transform/applications/P1/pytorch.cu +++ b/cub/benchmarks/bench/transform/applications/P1/pytorch.cu @@ -35,6 +35,8 @@ #include #include +#include +#include #include #include #include @@ -600,7 +602,7 @@ try d_a, n, [alpha] __device__(T a, T b, T c) -> T { - return pointwise_op_impl(a, b, c, alpha, std::multiplies()); + return pointwise_op_impl(a, b, c, alpha, cuda::std::multiplies()); }, s); @@ -782,7 +784,7 @@ try [] __device__(T in_) -> T { using opmath_t = opmath_type; const opmath_t in = in_; - const auto min = std::min(opmath_t(0), in); + const auto min = cuda::std::min(opmath_t(0), in); const auto z = std::exp(-std::abs(in)); return min - std::log1p(z); }, @@ -921,7 +923,7 @@ try [zero, one_sixth, three, six] __device__(T self_val) -> T { using opmath_t = opmath_type; opmath_t x = static_cast(self_val); - return x * std::min(std::max(x + three, zero), six) * one_sixth; + return x * cuda::std::min(cuda::std::max(x + three, zero), six) * one_sixth; }, s); @@ -943,7 +945,7 @@ try [zero, one_sixth, three, six] __device__(T self_val) -> T { using opmath_t = opmath_type; opmath_t x = static_cast(self_val); - return std::min(std::max(x + three, zero), six) * one_sixth; + return cuda::std::min(cuda::std::max(x + three, zero), six) * one_sixth; }, s); From 704c500985745a53c2efe0ee1f985aa810242fc8 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 10 Jul 2026 15:51:17 +0200 Subject: [PATCH 14/14] Use round_to_nearest_even on host and below SM80 --- .../bench/transform/applications/P1/bfloat16.h | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/cub/benchmarks/bench/transform/applications/P1/bfloat16.h b/cub/benchmarks/bench/transform/applications/P1/bfloat16.h index 275fcfe5dcc..880a9d6d87b 100644 --- a/cub/benchmarks/bench/transform/applications/P1/bfloat16.h +++ b/cub/benchmarks/bench/transform/applications/P1/bfloat16.h @@ -69,12 +69,12 @@ struct alignas(2) BFloat16 inline __host__ __device__ BFloat16::BFloat16(float value) { - NV_IF_TARGET(NV_IS_DEVICE, - ({ - __nv_bfloat16 tmp = __float2bfloat16(value); - x = *reinterpret_cast(&tmp); - }), - ({ x = bf16_detail::round_to_nearest_even(value); })); + NV_IF_ELSE_TARGET(NV_PROVIDES_SM_80, + ({ + __nv_bfloat16 tmp = __float2bfloat16(value); + x = *reinterpret_cast(&tmp); + }), + ({ x = bf16_detail::round_to_nearest_even(value); })); } inline __host__ __device__ BFloat16::operator float() const