From 89fac267789876d9bee9a746ba5c677a8de0f2b5 Mon Sep 17 00:00:00 2001 From: fbusato Date: Wed, 23 Sep 2026 16:53:45 -0700 Subject: [PATCH 1/2] extend large integer redux to signed types --- cub/cub/warp/specializations/warp_redux.cuh | 14 ++++++++------ cub/test/warp/catch2_test_warp_reduce.cu | 13 ++++++++----- 2 files changed, 16 insertions(+), 11 deletions(-) diff --git a/cub/cub/warp/specializations/warp_redux.cuh b/cub/cub/warp/specializations/warp_redux.cuh index 1a85b561cd45..e6dae82c382d 100644 --- a/cub/cub/warp/specializations/warp_redux.cuh +++ b/cub/cub/warp/specializations/warp_redux.cuh @@ -27,7 +27,7 @@ #include #include #include -#include +#include #include #include @@ -45,7 +45,7 @@ inline constexpr bool is_warp_redux_op_supported_sm80 = template > inline constexpr bool is_warp_redux_bitwise_large_supported = - ::cuda::std::is_unsigned_v && sizeof(T) > sizeof(unsigned) && is_cuda_std_bitwise_v; + ::cuda::std::is_integral_v && sizeof(T) > sizeof(unsigned) && is_cuda_std_bitwise_v; template > inline constexpr bool is_warp_redux_min_max_f32_supported = @@ -105,20 +105,22 @@ template warp_redux_bitwise_large(const T input, const ::cuda::std::uint32_t mask, ReductionOp reduction_op) { static_assert(is_warp_redux_bitwise_large_supported, "Reduction operator not supported"); + using unsigned_t = ::cuda::std::make_unsigned_t; // avoid signed shift UB constexpr int chunk_bits = 32; // 32 bits constexpr int num_chunks = sizeof(T) / sizeof(unsigned); const auto generalized_op = cub::detail::generalize_operator(reduction_op); // map bit_and to bit_and<> + const auto uinput = static_cast(input); // do not use bit_cast/memcpy to avoid potential performance issues - T output{}; + unsigned_t output{}; _CCCL_PRAGMA_UNROLL_FULL() for (int chunk = 0; chunk < num_chunks; ++chunk) { const int shift = chunk * chunk_bits; - const auto value = static_cast(input >> shift); + const auto value = static_cast(uinput >> shift); const auto reduced = cub::detail::warp_redux_sm80(value, mask, generalized_op); - output |= static_cast(reduced) << shift; + output |= static_cast(reduced) << shift; } - return output; + return static_cast(output); } //---------------------------------------------------------------------------------------------------------------------- diff --git a/cub/test/warp/catch2_test_warp_reduce.cu b/cub/test/warp/catch2_test_warp_reduce.cu index a3ffdd01802b..374fbddc76f8 100644 --- a/cub/test/warp/catch2_test_warp_reduce.cu +++ b/cub/test/warp/catch2_test_warp_reduce.cu @@ -191,10 +191,13 @@ using predefined_op_list = c2h::type_list, cuda::maximum<>, cu using predefined_min_max_op_list = c2h::type_list, cuda::minimum<>>; // clang-format off -using unsigned_type_list = c2h::type_list< +using bitwise_type_list = c2h::type_list< + int32_t, uint32_t, + int64_t, uint64_t #if TEST_INT128() + , __int128_t , __uint128_t #endif // TEST_INT128() >; @@ -325,10 +328,10 @@ CUB_TEST("WarpReduce::Max/Min, floating-point redux types", } } -CUB_TEST("WarpReduce::Reduce, unsigned bitwise types", +CUB_TEST("WarpReduce::Reduce, bitwise types", "[reduce][warp][predefined_op][redux]", CUB_SMALL, - unsigned_type_list, + bitwise_type_list, bitwise_op_list, logical_warp_threads) { @@ -392,10 +395,10 @@ CUB_TEST("WarpReduce::Sum/Max/Min Partial", verify_results(h_out, d_out); } -CUB_TEST("WarpReduce::Reduce, unsigned bitwise types, partial", +CUB_TEST("WarpReduce::Reduce, bitwise types, partial", "[reduce][warp][predefined_op][redux][partial]", CUB_SMALL, - unsigned_type_list, + bitwise_type_list, bitwise_op_list, logical_warp_threads) { From ed0792c3879967c77c143e5211a651a62140831f Mon Sep 17 00:00:00 2001 From: fbusato Date: Wed, 30 Sep 2026 18:22:38 -0700 Subject: [PATCH 2/2] workaround for GCC7 --- cub/test/warp/catch2_test_warp_reduce.cu | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/cub/test/warp/catch2_test_warp_reduce.cu b/cub/test/warp/catch2_test_warp_reduce.cu index 374fbddc76f8..35f261d88adc 100644 --- a/cub/test/warp/catch2_test_warp_reduce.cu +++ b/cub/test/warp/catch2_test_warp_reduce.cu @@ -328,6 +328,7 @@ CUB_TEST("WarpReduce::Max/Min, floating-point redux types", } } +#if _CCCL_COMPILER(GCC, >=, 8) // gcc 7 internal compiler error in test code only CUB_TEST("WarpReduce::Reduce, bitwise types", "[reduce][warp][predefined_op][redux]", CUB_SMALL, @@ -351,6 +352,8 @@ CUB_TEST("WarpReduce::Reduce, bitwise types", verify_results_exact(h_out, d_out); } +#endif // _CCCL_COMPILER(GCC, >=, 8) + CUB_TEST("WarpReduce::CustomSum", "[reduce][warp][generic][full]", CUB_SMALL, full_type_list, logical_warp_threads) { using T = c2h::get<0, TestType>; @@ -395,6 +398,8 @@ CUB_TEST("WarpReduce::Sum/Max/Min Partial", verify_results(h_out, d_out); } +#if _CCCL_COMPILER(GCC, >=, 8) // gcc 7 internal compiler error in test code only + CUB_TEST("WarpReduce::Reduce, bitwise types, partial", "[reduce][warp][predefined_op][redux][partial]", CUB_SMALL, @@ -419,6 +424,8 @@ CUB_TEST("WarpReduce::Reduce, bitwise types, partial", verify_results_exact(h_out, d_out); } +#endif // _CCCL_COMPILER(GCC, >=, 8) + CUB_TEST("WarpReduce::Sum", "[reduce][warp][generic][partial]", CUB_SMALL, full_type_list, logical_warp_threads) { using T = c2h::get<0, TestType>;