DO NOT MERGE: feat(hipcub, rocprim, rocthrust): cccl 3.x parity

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,35 @@
 
 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.2.0 (unreleased)
+
+### Added
+
+* Feature parity with CCCL/CUB 3.0.0.
+* 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
+
+* hipCUB compatibility with PyTorch v2.9 and v2.10 has been removed in this release.  Please use PyTorch v2.11 or later.
+* 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.7.0 for ROCm 10.1
 
 ### Changed

projects/hipcub/CMakeLists.txt

@@ -35,13 +35,13 @@ endif()
 # Update these variables at release time
 #
 # Set the library version
-set(VERSION_STRING "4.7.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(MIN_ROCPRIM_PACKAGE_VERSION "5.0.0" CACHE STRING "Minimum version of rocPRIM to search for when ROCPRIM_FETCH_METHOD is set to PACKAGE.")
 # Set download branch for dependency rocPRIM
 set(ROCM_DEP_RELEASE_BRANCH "release/therock-10.0" CACHE STRING "Download branch for ROCm dependencies")
 # --------------------------------------
@@ -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

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
@@ -65,7 +65,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];
@@ -111,7 +111,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];
@@ -159,7 +159,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];
@@ -205,7 +205,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];
@@ -385,7 +385,7 @@ void add_benchmarks(primbench::executor& executor)
     BENCHMARK_TYPE(int64_t, 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>)
     {
         BENCHMARK_TYPE(int, 256, true);
         BENCHMARK_TYPE(float, 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
@@ -25,17 +25,16 @@
 #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
 
 constexpr unsigned int Trials = 100;
 
-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;
     }
 };
 
@@ -60,18 +59,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));
-            }
-            else
+                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++)
@@ -98,18 +96,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));
-            }
-            else
+                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++)
@@ -136,23 +133,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
@@ -49,7 +49,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;
@@ -76,7 +76,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;
@@ -103,7 +103,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;
@@ -130,7 +130,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;
@@ -159,7 +159,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;
@@ -188,7 +188,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
@@ -60,13 +60,13 @@ struct histogram
         __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
@@ -46,7 +46,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;
@@ -72,7 +72,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;

projects/hipcub/benchmark/benchmark_block_radix_rank.cpp

@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2022-2024 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2022-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
@@ -67,15 +67,15 @@ void rank_kernel(const T* keys_input, int* ranks_output)
                                Descending,
                                BenchmarkKind == RadixRankAlgorithm::RADIX_RANK_MEMOIZE>>;
 
-#pragma unroll
+    _CCCL_PRAGMA_UNROLL_FULL()
     for(unsigned int key = 0; key < ItemsPerThread; key++)
     {
         unsigned_keys[key] = KeyTraits::TwiddleIn(unsigned_keys[key]);
     }
 
     int ranks[ItemsPerThread];
 
-#pragma nounroll
+    _CCCL_PRAGMA_NOUNROLL()
     for(unsigned int trial = 0; trial < Trials; trial++)
     {
         __shared__ typename RankType::TempStorage storage;