diff --git a/cub/cub/warp/specializations/warp_redux.cuh b/cub/cub/warp/specializations/warp_redux.cuh index 3fd1458e1031..eb49f6c08286 100644 --- a/cub/cub/warp/specializations/warp_redux.cuh +++ b/cub/cub/warp/specializations/warp_redux.cuh @@ -28,7 +28,7 @@ #include #include #include -#include +#include #include #include @@ -46,7 +46,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_plus_64bit_supported = @@ -198,20 +198,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 a1b30606190e..07d1978e30e0 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() >; @@ -357,10 +360,11 @@ CUB_TEST("WarpReduce::Max/Min, floating-point redux types", } } -CUB_TEST("WarpReduce::Reduce, unsigned bitwise 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, - unsigned_type_list, + bitwise_type_list, bitwise_op_list, logical_warp_threads) { @@ -380,6 +384,8 @@ CUB_TEST("WarpReduce::Reduce, unsigned 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>; @@ -424,10 +430,12 @@ CUB_TEST("WarpReduce::Sum/Max/Min Partial", verify_results(h_out, d_out); } -CUB_TEST("WarpReduce::Reduce, unsigned bitwise types, partial", +#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, - unsigned_type_list, + bitwise_type_list, bitwise_op_list, logical_warp_threads) { @@ -448,6 +456,8 @@ CUB_TEST("WarpReduce::Reduce, unsigned 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>;