diff --git a/cub/test/catch2_test_device_reduce.cuh b/cub/test/catch2_test_device_reduce.cuh index 9a3d2d887de7..d90b47c57321 100644 --- a/cub/test/catch2_test_device_reduce.cuh +++ b/cub/test/catch2_test_device_reduce.cuh @@ -17,143 +17,12 @@ #include #include +#include "catch2_test_numeric_helpers.cuh" #include #include #include #include -// Comparing results computed on CPU and GPU for extended floating point types is impossible. -// For instance, when used with a constant iterator of two, the accumulator in sequential reference -// computation (CPU) bumps into the 4096 limits, which will never change (`4096 + 2 = 4096`). -// Meanwhile, per-thread aggregates (`2 * 16 = 32`) are accumulated within and among thread blocks, -// yielding `inf` as a result. No reasonable epsilon can be selected to compare `inf` with `4096`. -// To make `__half` and `__nv_bfloat16` arithmetic associative, the function object below raises -// extended floating points to the area of unsigned short integers. This allows us to test large -// inputs with few code-path differences in device algorithms. -struct ExtendedFloatSum -{ - template - __host__ __device__ T operator()(T a, T b) const - { - T result{}; - result.__x = a.raw() + b.raw(); - return result; - } - -#if TEST_HALF_T() - __host__ __device__ __half operator()(__half a, __half b) const - { - uint16_t result = this->operator()(half_t{a}, half_t(b)).raw(); - return reinterpret_cast<__half&>(result); - } -#endif // TEST_HALF_T() - -#if TEST_BF_T() - __device__ __nv_bfloat16 operator()(__nv_bfloat16 a, __nv_bfloat16 b) const - { - uint16_t result = this->operator()(bfloat16_t{a}, bfloat16_t(b)).raw(); - return reinterpret_cast<__nv_bfloat16&>(result); - } -#endif // TEST_BF_T() -}; - -template -inline It unwrap_it(It it) -{ - return it; -} - -#if TEST_HALF_T() -inline __half* unwrap_it(half_t* it) -{ - return reinterpret_cast<__half*>(it); -} - -template -inline cuda::constant_iterator<__half, OffsetT> unwrap_it(cuda::constant_iterator it) -{ - half_t wrapped_val = *it; - __half val = wrapped_val.operator __half(); - return cuda::constant_iterator<__half, OffsetT>(val); -} -#endif // TEST_HALF_T() - -#if TEST_BF_T() -inline __nv_bfloat16* unwrap_it(bfloat16_t* it) -{ - return reinterpret_cast<__nv_bfloat16*>(it); -} - -template -cuda::constant_iterator<__nv_bfloat16, OffsetT> inline unwrap_it(cuda::constant_iterator it) -{ - bfloat16_t wrapped_val = *it; - __nv_bfloat16 val = wrapped_val.operator __nv_bfloat16(); - return cuda::constant_iterator<__nv_bfloat16, OffsetT>(val); -} -#endif // TEST_BF_T() - -template -using unwrap_value_t = std::remove_reference_t()))>; - -template ()))> -std::integral_constant> // - inline reference_extended_fp(WrappedItT) -{ - return {}; -} - -inline constexpr ExtendedFloatSum unwrap_op(std::true_type /* extended float */, ::cuda::std::plus<>) // -{ - return {}; -} - -template -inline constexpr OpT unwrap_op(std::integral_constant /* base case */, OpT op) -{ - return op; -} - -/** - * @brief Initializes the given item type with a constant non-zero value. - */ -template -inline void init_default_constant(T& val, int element_val = 2) -{ - val = T{static_cast(element_val)}; -} - -template