Skip to content

Fix [QA Agent] Cooperative reduction - Warp #11645

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
fbusato merged 3 commits into from
Oct 2, 2026
Merged

Fix [QA Agent] Cooperative reduction - Warp #11645

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
89fac26
extend large integer redux to signed types
fbusato Sep 23, 2026
ed0792c
workaround for GCC7
fbusato Oct 1, 2026
5a5e22f
Merge branch 'main' into fix-qa-agent-reduce
fbusato Oct 1, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
14 changes: 8 additions & 6 deletions cub/cub/warp/specializations/warp_redux.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,7 +28,7 @@
#include <cuda/std/__type_traits/is_integral.h>
#include <cuda/std/__type_traits/is_same.h>
#include <cuda/std/__type_traits/is_signed.h>
#include <cuda/std/__type_traits/is_unsigned.h>
#include <cuda/std/__type_traits/make_unsigned.h>
#include <cuda/std/__type_traits/remove_cvref.h>
#include <cuda/std/cstdint>

Expand All @@ -46,7 +46,7 @@ inline constexpr bool is_warp_redux_op_supported_sm80 =

template <typename Op, typename T, typename ReduceOp = ::cuda::std::remove_cvref_t<Op>>
inline constexpr bool is_warp_redux_bitwise_large_supported =
::cuda::std::is_unsigned_v<T> && sizeof(T) > sizeof(unsigned) && is_cuda_std_bitwise_v<ReduceOp, T>;
::cuda::std::is_integral_v<T> && sizeof(T) > sizeof(unsigned) && is_cuda_std_bitwise_v<ReduceOp, T>;

template <typename Op, typename T, typename ReduceOp = ::cuda::std::remove_cvref_t<Op>>
inline constexpr bool is_warp_redux_plus_64bit_supported =
Expand Down Expand Up @@ -198,20 +198,22 @@ template <typename T, typename ReductionOp>
warp_redux_bitwise_large(const T input, const ::cuda::std::uint32_t mask, ReductionOp reduction_op)
{
static_assert(is_warp_redux_bitwise_large_supported<ReductionOp, T>, "Reduction operator not supported");
using unsigned_t = ::cuda::std::make_unsigned_t<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<uint64_t> to bit_and<>
const auto uinput = static_cast<unsigned_t>(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<unsigned>(input >> shift);
const auto value = static_cast<unsigned>(uinput >> shift);
const auto reduced = cub::detail::warp_redux_sm80(value, mask, generalized_op);
output |= static_cast<T>(reduced) << shift;
output |= static_cast<unsigned_t>(reduced) << shift;
}
return output;
return static_cast<T>(output);
}

//----------------------------------------------------------------------------------------------------------------------
Expand Down
20 changes: 15 additions & 5 deletions cub/test/warp/catch2_test_warp_reduce.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -191,10 +191,13 @@ using predefined_op_list = c2h::type_list<cuda::std::plus<>, cuda::maximum<>, cu
using predefined_min_max_op_list = c2h::type_list<cuda::maximum<>, 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()
>;
Expand Down Expand Up @@ -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

@coderabbitai coderabbitai Bot Oct 1, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🎯 Functional Correctness | 🟡 Minor | ⚡ Quick win

🔎 Supported by static analysis

🏁 Script executed:

rg -n "_CCCL_COMPILER|GCC 7|catch2_test_warp_reduce|warp_reduce" libcudacxx/include/cuda/std/__cccl/compiler.h cub/test .github cmake
sed -n '350,470p' cub/test/warp/catch2_test_warp_reduce.cu

Repository: NVIDIA/cccl

Length of output: 35949

🏁 Script executed:

set -eu
printf '%s\n' '--- compiler macro semantics ---'
sed -n '1,115p' libcudacxx/include/cuda/std/__cccl/compiler.h
printf '%s\n' '--- warp test configuration references ---'
rg -n -C 3 'catch2_test_warp_reduce|warp_reduce.*(test|target)|CUB.*(GCC|Clang|MSVC|NVHPC)|NVHPC|Clang|MSVC|GCC 7|compiler' cub/test/CMakeLists.txt cub/test/warp .github cmake --glob '*.cmake' --glob '*.yml' --glob '*.yaml' --glob 'CMakeLists.txt' | head -n 300
printf '%s\n' '--- relevant test file setup and type lists ---'
sed -n '1,180p' cub/test/warp/catch2_test_warp_reduce.cu

Repository: NVIDIA/cccl

Length of output: 33792

🏁 Script executed:

set -eu
printf '%s\n' '--- CUB test target and source registration ---'
sed -n '1,230p' cub/test/CMakeLists.txt
printf '%s\n' '--- warp test references outside the first scope ---'
rg -n -C 3 'warp_reduce|CUB_TEST|CMAKE_CXX_COMPILER_ID|CMAKE_CUDA_COMPILER_ID|host:' .github cmake ci cub --glob '*.yml' --glob '*.yaml' --glob '*.cmake' --glob 'CMakeLists.txt' --glob '*.json' | head -n 500

Repository: NVIDIA/cccl

Length of output: 21357

important: Exclude only GCC versions below 8.

_CCCL_COMPILER(GCC, >=, 8) evaluates false for GCC below 8 and for inactive compilers. This removes both bitwise reduction tests for Clang, MSVC, and NVHPC, although the generic CUB test configuration builds this test source for those compilers.

Use !_CCCL_COMPILER(GCC, <, 8) at both guards and update the matching #endif annotations. The negated comparison evaluates true for inactive compilers.

Suggested fix 
-#if _CCCL_COMPILER(GCC, >=, 8) // gcc 7 internal compiler error in test code only
+#if !_CCCL_COMPILER(GCC, <, 8) // gcc 7 internal compiler error in test code only
...
-#endif // _CCCL_COMPILER(GCC, >=, 8)
+#endif // !_CCCL_COMPILER(GCC, <, 8)
...
-#if _CCCL_COMPILER(GCC, >=, 8) // gcc 7 internal compiler error in test code only
+#if !_CCCL_COMPILER(GCC, <, 8) // gcc 7 internal compiler error in test code only
...
-#endif // _CCCL_COMPILER(GCC, >=, 8)
+#endif // !_CCCL_COMPILER(GCC, <, 8)
📝 Committable suggestion

‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.

Suggested change
#if _CCCL_COMPILER(GCC, >=, 8) // gcc 7 internal compiler error in test code only
#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)
{
Expand All @@ -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>;
Expand Down Expand Up @@ -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)
{
Expand All @@ -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>;
Expand Down
Loading