feat(hipcub): Add CCCL 3.0.x support (copy)

projects/hipcub/.clang-format

@@ -167,4 +167,10 @@ Macros:
     - HIPCUB_DETAIL_DEPRECATED_DEBUG_SYNCHRONOUS=[[DETAIL_DEPRECATED_DEBUG_SYNCHRONOUS___]]
 BreakAfterAttributes: Always
 
+WhitespaceSensitiveMacros: [
+  'HIPCUB_HAS_INCLUDE',
+  '_HIPCUB_LIBCXX_INCLUDE',
+  '_HIPCUB_STD_INCLUDE'
+]
+
 ---

projects/hipcub/CHANGELOG.md

@@ -2,6 +2,33 @@
 
 Full documentation for hipCUB is available at [https://rocm.docs.amd.com/projects/hipCUB/en/latest/](https://rocm.docs.amd.com/projects/hipCUB/en/latest/).
 
+## hipCUB-5.0.0 for ROCm 10.0
+
+### Added
+
+* Added `::hip::std` support.
+
+### Changed
+
+* Changed `CCCL_MINIMUM_VERSION` to `3.0.0` to align with CUB.
+* Add support for large num_items `DeviceMerge` and `DeviceSegmentedSort`.
+* Replace `#pragma unroll` by `_CCCL_PRAGMA_UNROLL_FULL()` and `_CCCL_PRAGMA_NOUNROLL()` by `_CCCL_PRAGMA_NOUNROLL()`.
+* Add `_CCCL_SORT_MAYBE_UNROLL()` in block merge sort and thread sort.
+* Update `WarpExchange` template parameters for CUB compatibility.
+
+### Removed
+
+* Removed `hipcub::BaseTraits::CATEGORY`, `hipcub::BaseTraits::nullptr_TYPE` and `hipcub::BaseTraits::PRIMITIVE`.
+* Removed  `ConstantInputIterator`, `CountingInputIterator`, `DiscardOutputIterator` and `TransformInputIterator` which were deprecated in hipCUB-4.1.0.
+* Removed `DeviceSpmv`, which was removed from CUB after CCCL's 2.8.0 release. Use `hipSPARSE` or `rocSPARSE` libraries instead.
+* Removed `GridBarrier`.
+* Removed `HIPCUB_MIN`, `HIPCUB_MAX`, `HIPCUB_QUOTIENT_FLOOR`, `HIPCUB_QUOTIENT_CEILING`, `HIPCUB_ROUND_UP_NEAREST` and `HIPCUB_ROUND_DOWN_NEAREST` which were deprecated in hipCUB-4.1.0.
+* Removed `LEGACY_PTX_ARCH`.
+* Removed `hipcub:max` and `hipcub:min`, which were deprecated. Use `hip::std::max` and `hip::std::min` instead.
+* Deprecated `hipcub::Swap`, use `rocprim::swap` instead.
+* Deprecated `HIPCUB_IS_INT128_ENABLED`, use `_CCCL_HAS_INT128()` instead.
+* Deprecated `hipcub::Equality`, `hipcub::Inequality`, `hipcub::InequalityWrapper`, `hipcub::Sum`, `hipcub::Difference`, `hipcub::Division`, `hipcub::Max` and `hipcub::Min` operators. Use `hip::std::equal_to`, `hip::std::not_equal_to`, `hip::std::plus`, `hip::std::minus`, `hip::std::divides`, `hip::maximum` and `hip:minimum` operators instead.
+
 ## hipCUB 4.5.0 for ROCm 7.14
 
 ### Added

projects/hipcub/CMakeLists.txt

@@ -35,11 +35,11 @@ endif()
 # Update these variables at release time
 #
 # Set the library version
-set(VERSION_STRING "4.5.0")
+set(VERSION_STRING "5.0.0")
 # Set the CCCL-compatible version.
-set(HIPCUB_CCCL_VERSION_MAJOR 2)
-set(HIPCUB_CCCL_VERSION_MINOR 8)
-set(HIPCUB_CCCL_VERSION_PATCH 2)
+set(HIPCUB_CCCL_VERSION_MAJOR 3)
+set(HIPCUB_CCCL_VERSION_MINOR 0)
+set(HIPCUB_CCCL_VERSION_PATCH 3)
 # Set the minimum required rocPRIM version
 set(MIN_ROCPRIM_PACKAGE_VERSION "4.1.0" CACHE STRING "Minimum version of rocPRIM to search for when ROCPRIM_FETCH_METHOD is set to PACKAGE.")
 # Set download branch for dependency rocPRIM
@@ -59,10 +59,8 @@ endif()
 set(CMAKE_CXX_STANDARD_REQUIRED ON)
 set(CMAKE_CXX_EXTENSIONS OFF)
 
-if (CMAKE_CXX_STANDARD EQUAL 14)
-  message(WARNING "C++14 will be deprecated in the next major release")
-elseif(NOT CMAKE_CXX_STANDARD EQUAL 17)
-  message(FATAL_ERROR "Only C++14 and C++17 are supported")
+if(NOT CMAKE_CXX_STANDARD EQUAL 17)
+  message(FATAL_ERROR "Only C++17 are supported")
 endif()
 
 # Set HIP flags

projects/hipcub/benchmark/CMakeLists.txt

@@ -1,6 +1,6 @@
 # MIT License
 #
-# Copyright (c) 2020-2025 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (c) 2020-2026 Advanced Micro Devices, Inc. All rights reserved.
 #
 # Permission is hereby granted, free of charge, to any person obtaining a copy
 # of this software and associated documentation files (the "Software"), to deal
@@ -35,6 +35,11 @@ function(add_hipcub_benchmark BENCHMARK_SOURCE)
       benchmark::benchmark
       hipcub
   )
+
+  if (WIN32)
+    target_compile_definitions(${BENCHMARK_TARGET} PRIVATE BENCHMARK_STATIC_DEFINE)
+  endif()
+
   if((HIP_COMPILER STREQUAL "nvcc"))
     set_property(TARGET ${BENCHMARK_TARGET} PROPERTY CUDA_STANDARD 17)
     set_source_files_properties(${BENCHMARK_SOURCE} PROPERTIES LANGUAGE CUDA)
@@ -100,7 +105,6 @@ add_hipcub_benchmark(benchmark_device_segmented_sort.cpp)
 add_hipcub_benchmark(benchmark_device_segmented_radix_sort.cpp)
 add_hipcub_benchmark(benchmark_device_segmented_reduce.cpp)
 add_hipcub_benchmark(benchmark_device_select.cpp)
-add_hipcub_benchmark(benchmark_device_spmv.cpp)
 add_hipcub_benchmark(benchmark_warp_exchange.cpp)
 add_hipcub_benchmark(benchmark_warp_load.cpp)
 add_hipcub_benchmark(benchmark_warp_reduce.cpp)

projects/hipcub/benchmark/benchmark_block_adjacent_difference.cpp

@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2020-2022 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2020-2026 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -63,7 +63,7 @@ struct subtract_left
 
         hipcub::BlockAdjacentDifference<T, BlockSize> adjacent_difference;
 
-#pragma nounroll
+        _CCCL_PRAGMA_NOUNROLL()
         for(unsigned int trial = 0; trial < trials; trial++)
         {
             T output[ItemsPerThread];
@@ -106,7 +106,7 @@ struct subtract_left_partial_tile
         // Try to evenly distribute the length of tile_sizes between all the trials
         const auto tile_size_diff = (BlockSize * ItemsPerThread) / trials + 1;
 
-#pragma nounroll
+        _CCCL_PRAGMA_NOUNROLL()
         for(unsigned int trial = 0; trial < trials; trial++)
         {
             T output[ItemsPerThread];
@@ -150,7 +150,7 @@ struct subtract_right
 
         hipcub::BlockAdjacentDifference<T, BlockSize> adjacent_difference;
 
-#pragma nounroll
+        _CCCL_PRAGMA_NOUNROLL()
         for(unsigned int trial = 0; trial < trials; trial++)
         {
             T output[ItemsPerThread];
@@ -193,7 +193,7 @@ struct subtract_right_partial_tile
         // Try to evenly distribute the length of tile_sizes between all the trials
         const auto tile_size_diff = (BlockSize * ItemsPerThread) / trials + 1;
 
-#pragma nounroll
+        _CCCL_PRAGMA_NOUNROLL()
         for(unsigned int trial = 0; trial < trials; trial++)
         {
             T output[ItemsPerThread];
@@ -221,8 +221,8 @@ template<class Benchmark,
          bool         WithTile,
          unsigned int Trials = 100>
 auto run_benchmark(benchmark::State& state, hipStream_t stream, size_t N)
-    -> std::enable_if_t<!std::is_same<Benchmark, subtract_left_partial_tile>::value
-                        && !std::is_same<Benchmark, subtract_right_partial_tile>::value>
+    -> std::enable_if_t<!std::is_same_v<Benchmark, subtract_left_partial_tile>
+                        && !std::is_same_v<Benchmark, subtract_right_partial_tile>>
 {
     constexpr auto items_per_block = BlockSize * ItemsPerThread;
     const auto     num_blocks      = (N + items_per_block - 1) / items_per_block;
@@ -271,8 +271,8 @@ template<class Benchmark,
          bool         WithTile,
          unsigned int Trials = 100>
 auto run_benchmark(benchmark::State& state, hipStream_t stream, size_t N)
-    -> std::enable_if_t<std::is_same<Benchmark, subtract_left_partial_tile>::value
-                        || std::is_same<Benchmark, subtract_right_partial_tile>::value>
+    -> std::enable_if_t<std::is_same_v<Benchmark, subtract_left_partial_tile>
+                        || std::is_same_v<Benchmark, subtract_right_partial_tile>>
 {
     constexpr auto items_per_block = BlockSize * ItemsPerThread;
     const auto     num_blocks      = (N + items_per_block - 1) / items_per_block;
@@ -352,7 +352,7 @@ void add_benchmarks(const std::string&                            name,
                                                        BENCHMARK_TYPE(long long, 256, false),
                                                        BENCHMARK_TYPE(double, 256, false)};
 
-    if(!std::is_same<Benchmark, subtract_right_partial_tile>::value)
+    if(!std::is_same_v<Benchmark, subtract_right_partial_tile>)
     {
         bs.insert(bs.end(),
                   {BENCHMARK_TYPE(int, 256, true),

projects/hipcub/benchmark/benchmark_block_discontinuity.cpp

@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2020 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2020-2026 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -24,21 +24,20 @@
 #include <hipcub/block/block_discontinuity.hpp>
 #include <hipcub/block/block_load.hpp>
 #include <hipcub/block/block_store.hpp>
-#include <hipcub/thread/thread_operators.hpp> //to use hipcub::Equality
 
 #include "common_benchmark_header.hpp"
 
 #ifndef DEFAULT_N
 const size_t DEFAULT_N = 1024 * 1024 * 128;
 #endif
 
-template<class T>
-struct custom_flag_op1
+struct equal
 {
-    HIPCUB_HOST_DEVICE
-    bool operator()(const T& a, const T& b) const
+    template<class A, class B>
+        HIPCUB_HOST_DEVICE
+    inline constexpr auto operator()(const A& a, const B& b) const
     {
-        return (a == b);
+        return a == b;
     }
 };
 
@@ -68,17 +67,17 @@ struct flag_heads
         T input[ItemsPerThread];
         hipcub::LoadDirectStriped<BlockSize>(lid, d_input + block_offset, input);
 
-#pragma nounroll
+        _CCCL_PRAGMA_NOUNROLL()
         for(unsigned int trial = 0; trial < Trials; trial++)
         {
             hipcub::BlockDiscontinuity<T, BlockSize> bdiscontinuity;
             bool                                     head_flags[ItemsPerThread];
             if(WithTile)
             {
-                bdiscontinuity.FlagHeads(head_flags, input, hipcub::Equality(), T(123));
+                bdiscontinuity.FlagHeads(head_flags, input, equal(), T(123));
             } else
             {
-                bdiscontinuity.FlagHeads(head_flags, input, hipcub::Equality());
+                bdiscontinuity.FlagHeads(head_flags, input, equal());
             }
 
             for(unsigned int i = 0; i < ItemsPerThread; i++)
@@ -106,17 +105,17 @@ struct flag_tails
         T input[ItemsPerThread];
         hipcub::LoadDirectStriped<BlockSize>(lid, d_input + block_offset, input);
 
-#pragma nounroll
+        _CCCL_PRAGMA_NOUNROLL()
         for(unsigned int trial = 0; trial < Trials; trial++)
         {
             hipcub::BlockDiscontinuity<T, BlockSize> bdiscontinuity;
             bool                                     tail_flags[ItemsPerThread];
             if(WithTile)
             {
-                bdiscontinuity.FlagTails(tail_flags, input, hipcub::Equality(), T(123));
+                bdiscontinuity.FlagTails(tail_flags, input, equal(), T(123));
             } else
             {
-                bdiscontinuity.FlagTails(tail_flags, input, hipcub::Equality());
+                bdiscontinuity.FlagTails(tail_flags, input, equal());
             }
 
             for(unsigned int i = 0; i < ItemsPerThread; i++)
@@ -144,23 +143,20 @@ struct flag_heads_and_tails
         T input[ItemsPerThread];
         hipcub::LoadDirectStriped<BlockSize>(lid, d_input + block_offset, input);
 
-#pragma nounroll
+        _CCCL_PRAGMA_NOUNROLL()
         for(unsigned int trial = 0; trial < Trials; trial++)
         {
             hipcub::BlockDiscontinuity<T, BlockSize> bdiscontinuity;
             bool                                     head_flags[ItemsPerThread];
             bool                                     tail_flags[ItemsPerThread];
             if(WithTile)
             {
-                bdiscontinuity.FlagHeadsAndTails(head_flags,
-                                                 T(123),
-                                                 tail_flags,
-                                                 T(234),
-                                                 input,
-                                                 hipcub::Equality());
-            } else
+                bdiscontinuity
+                    .FlagHeadsAndTails(head_flags, T(123), tail_flags, T(234), input, equal());
+            }
+            else
             {
-                bdiscontinuity.FlagHeadsAndTails(head_flags, tail_flags, input, hipcub::Equality());
+                bdiscontinuity.FlagHeadsAndTails(head_flags, tail_flags, input, equal());
             }
 
             for(unsigned int i = 0; i < ItemsPerThread; i++)

projects/hipcub/benchmark/benchmark_block_exchange.cpp

@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2020 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2020-2026 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -54,7 +54,7 @@ struct blocked_to_striped
         T input[ItemsPerThread];
         hipcub::LoadDirectBlocked(lid, d_input + block_offset, input);
 
-#pragma nounroll
+        _CCCL_PRAGMA_NOUNROLL()
         for(unsigned int trial = 0; trial < Trials; trial++)
         {
             hipcub::BlockExchange<T, BlockSize, ItemsPerThread> exchange;
@@ -78,7 +78,7 @@ struct striped_to_blocked
         T input[ItemsPerThread];
         hipcub::LoadDirectStriped<BlockSize>(lid, d_input + block_offset, input);
 
-#pragma nounroll
+        _CCCL_PRAGMA_NOUNROLL()
         for(unsigned int trial = 0; trial < Trials; trial++)
         {
             hipcub::BlockExchange<T, BlockSize, ItemsPerThread> exchange;
@@ -102,7 +102,7 @@ struct blocked_to_warp_striped
         T input[ItemsPerThread];
         hipcub::LoadDirectBlocked(lid, d_input + block_offset, input);
 
-#pragma nounroll
+        _CCCL_PRAGMA_NOUNROLL()
         for(unsigned int trial = 0; trial < Trials; trial++)
         {
             hipcub::BlockExchange<T, BlockSize, ItemsPerThread> exchange;
@@ -126,7 +126,7 @@ struct warp_striped_to_blocked
         T input[ItemsPerThread];
         hipcub::LoadDirectWarpStriped(lid, d_input + block_offset, input);
 
-#pragma nounroll
+        _CCCL_PRAGMA_NOUNROLL()
         for(unsigned int trial = 0; trial < Trials; trial++)
         {
             hipcub::BlockExchange<T, BlockSize, ItemsPerThread> exchange;
@@ -152,7 +152,7 @@ struct scatter_to_blocked
         hipcub::LoadDirectStriped<BlockSize>(lid, d_input + block_offset, input);
         hipcub::LoadDirectStriped<BlockSize>(lid, d_ranks + block_offset, ranks);
 
-#pragma nounroll
+        _CCCL_PRAGMA_NOUNROLL()
         for(unsigned int trial = 0; trial < Trials; trial++)
         {
             hipcub::BlockExchange<T, BlockSize, ItemsPerThread> exchange;
@@ -178,7 +178,7 @@ struct scatter_to_striped
         hipcub::LoadDirectStriped<BlockSize>(lid, d_input + block_offset, input);
         hipcub::LoadDirectStriped<BlockSize>(lid, d_ranks + block_offset, ranks);
 
-#pragma nounroll
+        _CCCL_PRAGMA_NOUNROLL()
         for(unsigned int trial = 0; trial < Trials; trial++)
         {
             hipcub::BlockExchange<T, BlockSize, ItemsPerThread> exchange;

projects/hipcub/benchmark/benchmark_block_histogram.cpp

@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2020 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2020-2026 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -64,13 +64,13 @@ struct histogram
         __shared__ T                                  histogram[BinSize];
         __shared__ typename bhistogram_t::TempStorage storage;
 
-#pragma nounroll
+        _CCCL_PRAGMA_NOUNROLL()
         for(unsigned int trial = 0; trial < Trials; trial++)
         {
             bhistogram_t(storage).Histogram(values, histogram);
         }
 
-#pragma unroll
+        _CCCL_PRAGMA_UNROLL_FULL()
         for(unsigned int offset = 0; offset < BinSize; offset += BlockSize)
         {
             if(offset + hipThreadIdx_x < BinSize)

projects/hipcub/benchmark/benchmark_block_merge_sort.cpp

@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2021-2024 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2021-2026 Advanced Micro Devices, Inc. All rights reserved.
 //
 // Permission is hereby granted, free of charge, to any person obtaining a copy
 // of this software and associated documentation files (the "Software"), to deal
@@ -52,7 +52,7 @@ void sort_keys_kernel(const T* input, T* output, CompareOp compare_op)
     T keys[ItemsPerThread];
     hipcub::LoadDirectStriped<BlockSize>(lid, input + block_offset, keys);
 
-#pragma nounroll
+    _CCCL_PRAGMA_NOUNROLL()
     for(unsigned int trial = 0; trial < Trials; trial++)
     {
         hipcub::BlockMergeSort<T, BlockSize, ItemsPerThread> sort;
@@ -82,7 +82,7 @@ void sort_pairs_kernel(const T* input, T* output, CompareOp compare_op)
         values[i] = keys[i] + T(1);
     }
 
-#pragma nounroll
+    _CCCL_PRAGMA_NOUNROLL()
     for(unsigned int trial = 0; trial < Trials; trial++)
     {
         hipcub::BlockMergeSort<T, BlockSize, ItemsPerThread, T> sort;