Beanavil · Oct 19, 2026 · Oct 28, 2026 · Oct 6, 2026 · Oct 6, 2026 · Oct 7, 2026
@@ -2,7 +2,7 @@
 
 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.1.0 for ROCm x.y.z (unreleased)
+## hipCUB-5.1.0 for ROCm 10.3.0
 
 ### Added
 
@@ -16,16 +16,34 @@ Full documentation for hipCUB is available at [https://rocm.docs.amd.com/project
 * Added env-based API for `hipcub::DeviceReduce::Reduce`.
 * Added env-based API for `hipcub::DeviceReduce::Sum`.
 * Feature parity with CCCL/CUB 3.1.4.
+* Added env-based API for `hipcub::DeviceReduce::(Arg)MinMax`.
+* Added env-based API for `hipcub::DeviceScan::ExclusiveScan` and `hipcub::DeviceScan::ExclusiveSum`.
+* Added new `hipcub::DeviceTopK::{Max,Min}Keys` and `hipcub::DeviceTopK::{Max,Min}Pairs`.
+* Added `hipcub::DeviceFor::ForEachInLayout`.
+* Added `hipcub::DeviceCopy::Copy`.
+* Added `hipcub::WarpScan::InclusiveScanPartial`, `hipcub::WarpScan::ExclusiveScanPartial` and `hipcub::WarpScan::ScanPartial`.
+* Added `hipcub::DeviceTransform::TransformIf`, `hipcub::DeviceTransform::Generate` and `hipcub::DeviceTransform::Fill`.
+* Added env-based API for `hipcub::DeviceTransform::*`.
+* Feature parity with CCCL/CUB 3.2.0.
+* Feature parity with CCCL/CUB 3.2.1.
 
 ### Changed
 
 * Replaced `_CCCL_PRAGMA_UNROLL_FULL()` and `_CCCL_PRAGMA_NOUNROLL()` by `HIPCUB_PRAGMA_UNROLL_FULL()` and `HIPCUB_PRAGMA_NOUNROLL()`.
 * Replaced `_CCCL_SORT_MAYBE_UNROLL()` by `HIPCUB_SORT_MAYBE_UNROLL()`.
-* Changed `CCCL_MINIMUM_VERSION` to `3.1.4` to align with CUB.
+* Changed `CCCL_MINIMUM_VERSION` to `3.2.1` to align with CUB.
+* `::hipcub::extents` now points to `::hip::std::extents` from libhipcxx for rocPRIM's backend.
+* Removed deprecated `hipcub::DivideAndRoundUp` in favor of `hip::ceil_div`.
+* Fixed handling of invalid inputs ($\leq 0$) in `hipcub::GridEvenShare::DispatchInit()`, which now sets all relevant members to 0 when any of the inputs is invalid.
 
 ### Removed
 
-* Removed the deprecated thread operators `Equality`, `Inequality`, `Sum`, `Difference`, `Division`, `Max` and `Min`.
+* Removed the deprecated thread operators `Equality`, `Sum`, `Difference`, `Division`, `Max` and `Min`.
+* Removed the `HIPCUB_TRAIT` macro in favor of the equivalent `_CCCL_TRAIT` from libhipcxx.
+* Deprecated `hipcub::PowerOfTwo`, use `hip::is_power_of_two(N)` instead.
+* Removed `WARP_SORT` template argument from `hipcub::BlockMergeSortStrategy`.
+* Removed the deprecated `hipcub::AliasTemporaries`.
+* Deprecated `hipcub::MergePathSearch`.
 
 ## hipCUB-5.0.0 for ROCm 10.2.0
 
@@ -47,6 +65,7 @@ Full documentation for hipCUB is available at [https://rocm.docs.amd.com/project
 * Fixed a build failure when libhipcxx is not on the include path. `radix_rank_sort_operations.hpp` included the libhipcxx `bit` header unconditionally, which does not exist in the standard library fallback.
 
 ### Removed
+
 * The hipCUB migration to cccl3 has resulted in the following removals and deprecations:
   * 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.

@@ -47,13 +47,15 @@ endif()
 # already shipped at 4.x, so aligning it down to the umbrella version would
 # be a backwards version move and would break downstream
 # find_package(hipcub 4.x) consumers.
-set(VERSION_STRING "5.0.0")
+set(VERSION_STRING "5.1.0")
 # Set the CCCL-compatible version.
 set(HIPCUB_CCCL_VERSION_MAJOR 3)
-set(HIPCUB_CCCL_VERSION_MINOR 1)
-set(HIPCUB_CCCL_VERSION_PATCH 4)
+set(HIPCUB_CCCL_VERSION_MINOR 2)
+set(HIPCUB_CCCL_VERSION_PATCH 1)
 # Set the minimum required rocPRIM version
-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(MIN_ROCPRIM_PACKAGE_VERSION "5.1.0" CACHE STRING "Minimum version of rocPRIM to search for when ROCPRIM_FETCH_METHOD is set to PACKAGE.")
+# Set the minimum required libhipcxx version
+set(MIN_LIBHIPCXX_PACKAGE_VERSION "3.4.0" CACHE STRING "Minimum version of libhipcxx to search for when LIBHIPCXX_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")
 # --------------------------------------

@@ -69,7 +69,7 @@ struct subtract_left
         for(unsigned int trial = 0; trial < trials; trial++)
         {
             T output[ItemsPerThread];
-            if(WithTile)
+            if constexpr(WithTile)
             {
                 adjacent_difference.SubtractLeft(input, output, minus<T>{}, T(123));
             }
@@ -116,7 +116,7 @@ struct subtract_left_partial_tile
         {
             T output[ItemsPerThread];
 
-            if(WithTile)
+            if constexpr(WithTile)
             {
                 adjacent_difference.SubtractLeftPartialTile(input,
                                                             output,
@@ -163,7 +163,7 @@ struct subtract_right
         for(unsigned int trial = 0; trial < trials; trial++)
         {
             T output[ItemsPerThread];
-            if(WithTile)
+            if constexpr(WithTile)
             {
                 adjacent_difference.SubtractRight(input, output, minus<T>{}, T(123));
             }

@@ -64,7 +64,7 @@ struct flag_heads
         {
             hipcub::BlockDiscontinuity<T, BlockSize> bdiscontinuity;
             bool                                     head_flags[ItemsPerThread];
-            if(WithTile)
+            if constexpr(WithTile)
             {
                 bdiscontinuity.FlagHeads(head_flags, input, equal(), T(123));
             } else
@@ -101,7 +101,7 @@ struct flag_tails
         {
             hipcub::BlockDiscontinuity<T, BlockSize> bdiscontinuity;
             bool                                     tail_flags[ItemsPerThread];
-            if(WithTile)
+            if constexpr(WithTile)
             {
                 bdiscontinuity.FlagTails(tail_flags, input, equal(), T(123));
             } else
@@ -139,7 +139,7 @@ struct flag_heads_and_tails
             hipcub::BlockDiscontinuity<T, BlockSize> bdiscontinuity;
             bool                                     head_flags[ItemsPerThread];
             bool                                     tail_flags[ItemsPerThread];
-            if(WithTile)
+            if constexpr(WithTile)
             {
                 bdiscontinuity
                     .FlagHeadsAndTails(head_flags, T(123), tail_flags, T(234), input, equal());

@@ -248,7 +248,7 @@ class block_radix_sort_benchmark : public primbench::benchmark_interface
                         <<<dim3(items / items_per_block), dim3(BlockSize), 0, stream>>>(d_input,
                                                                                         d_output);
                 }
-                else if(BenchmarkKind == benchmark_kinds::sort_pairs)
+                else if constexpr(BenchmarkKind == benchmark_kinds::sort_pairs)
                 {
                     sort_pairs_kernel<Helper, T, BlockSize, ItemsPerThread>
                         <<<dim3(items / items_per_block), dim3(BlockSize), 0, stream>>>(d_input,

@@ -31,7 +31,7 @@
 #include <hipcub/device/device_reduce.hpp>
 #include <hipcub/device/device_scan.hpp>
 
-#include _HIPCUB_STD_INCLUDE(functional)
+#include _HIPCUB_STD_INCLUDE(__functional/operations.h)
 
 template<bool Exclusive, class T, class BinaryFunction>
 auto run_device_scan(void*             temporary_storage,

@@ -54,7 +54,7 @@
     #define HIPCUB_WARP_THREADS_MACRO warp_threads
 #endif
 
-#include _HIPCUB_STD_INCLUDE(limits)
+#include _HIPCUB_STD_INCLUDE(__limits/numeric_limits.h)
 
 namespace benchmark_utils
 {

@@ -426,7 +426,7 @@ endif(USER_BUILD_TEST)
 
 # CUB (only for CUDA platform)
 if(HIP_COMPILER STREQUAL "nvcc")
-  set(CCCL_MINIMUM_VERSION 3.1.4)
+  set(CCCL_MINIMUM_VERSION 3.2.1)
   if(NOT DOWNLOAD_CUB)
     find_package(CCCL ${CCCL_MINIMUM_VERSION} CONFIG)
   endif()
@@ -515,7 +515,7 @@ else()
   message(STATUS "Searching for ${REPO_NAME} package")
 
   # Add default install location for WIN32 and non-WIN32 as hint
-  find_package(${REPO_NAME} CONFIG QUIET PATHS "${ROCM_ROOT}/lib/cmake/${REPO_NAME}")
+  find_package(${REPO_NAME} ${MIN_LIBHIPCXX_PACKAGE_VERSION} CONFIG QUIET PATHS "${ROCM_ROOT}/lib/cmake/${REPO_NAME}")
 
   if(NOT ${${REPO_NAME}_FOUND})
     message(STATUS "No existing ${REPO_NAME} package was found.")

@@ -102,7 +102,7 @@ Removed functions and classes
    * - ``hipcub::DiscardOutputIterator``
      - Use ``thrust::discard_iterator`` instead.
    * - ``hipcub::DivideAndRoundUp``
-     - Use ``hip::round_up`` instead.
+     - Use ``hip::ceil_div`` instead.
    * - ``hipcub::Division``
      - Use ``hip::std::divides`` instead.
    * - ``hipcub::Equality``

@@ -46,6 +46,9 @@
 
 #include "../example_utils.hpp"
 
+#include _HIPCUB_STD_INCLUDE(__algorithm/max.h)
+#include _HIPCUB_STD_INCLUDE(__algorithm/min.h)
+
 using namespace hipcub;
 
 

@@ -46,6 +46,9 @@
 
 #include "../example_utils.hpp"
 
+#include _HIPCUB_STD_INCLUDE(__algorithm/max.h)
+#include _HIPCUB_STD_INCLUDE(__algorithm/min.h)
+
 using namespace hipcub;
 
 

@@ -47,6 +47,9 @@
 
 #include "../example_utils.hpp"
 
+#include _HIPCUB_STD_INCLUDE(__algorithm/max.h)
+#include _HIPCUB_STD_INCLUDE(__algorithm/min.h)
+
 using namespace hipcub;
 
 

@@ -47,6 +47,9 @@
 
 #include "../example_utils.hpp"
 
+#include _HIPCUB_STD_INCLUDE(__algorithm/max.h)
+#include _HIPCUB_STD_INCLUDE(__algorithm/min.h)
+
 using namespace hipcub;
 
 

@@ -40,7 +40,7 @@
 #include <sstream>
 #include <vector>
 
-#include _HIPCUB_STD_INCLUDE(functional)
+#include _HIPCUB_STD_INCLUDE(__algorithm/max.h)
 
 #define AssertEquals(a, b) if ((a) != (b)) { std::cerr << "\n(" << __FILE__ << ": " << __LINE__ << ")\n"; exit(1);}
 

@@ -56,6 +56,30 @@ struct DeviceCopy
                                                                  num_buffers,
                                                                  stream));
     }
+
+    template<typename T_In,
+             typename Extents_In,
+             typename Layout_In,
+             typename Accessor_In,
+             typename T_Out,
+             typename Extents_Out,
+             typename Layout_Out,
+             typename Accessor_Out>
+    [[nodiscard]]
+    HIPCUB_RUNTIME_FUNCTION
+    static hipError_t
+        Copy(void*                                                             d_temp_storage,
+             size_t&                                                           temp_storage_bytes,
+             _HIPCUB_STD::mdspan<T_In, Extents_In, Layout_In, Accessor_In>     mdspan_in,
+             _HIPCUB_STD::mdspan<T_Out, Extents_Out, Layout_Out, Accessor_Out> mdspan_out,
+             hipStream_t                                                       stream = nullptr)
+    {
+        return hipCUDAErrorTohipError(::cub::DeviceCopy::Copy(d_temp_storage,
+                                                              temp_storage_bytes,
+                                                              mdspan_in,
+                                                              mdspan_out,
+                                                              stream));
+    }
 };
 
 END_HIPCUB_NAMESPACE

@@ -186,6 +186,33 @@ HIPCUB_RUNTIME_FUNCTION
     {
         return hipCUDAErrorTohipError(cub::DeviceFor::ForEachInExtents(extents, op, stream));
     }
+
+    HIPCUB_TEMPLATE(class LayoutMapping, typename OpType)
+    HIPCUB_REQUIRES(_HIPCUB_STD::__is_cuda_std_layout_left_or_right_mapping_v<LayoutMapping>)
+    [[nodiscard]]
+    HIPCUB_RUNTIME_FUNCTION
+    static hipError_t
+        ForEachInLayout(const LayoutMapping& layout_mapping, OpType op, cudaStream_t stream = {})
+    {
+        return hipCUDAErrorTohipError(cub::DeviceFor::ForEachInLayout(layout_mapping, op, stream));
+    }
+
+    HIPCUB_TEMPLATE(class LayoutMapping, typename OpType)
+    HIPCUB_REQUIRES(_HIPCUB_STD::__is_cuda_std_layout_left_or_right_mapping_v<LayoutMapping>)
+    [[nodiscard]]
+    HIPCUB_RUNTIME_FUNCTION
+    static hipError_t ForEachInLayout(void*                d_temp_storage,
+                                      size_t&              temp_storage_bytes,
+                                      const LayoutMapping& layout_mapping,
+                                      OpType               op,
+                                      cudaStream_t         stream = {})
+    {
+        return hipCUDAErrorTohipError(cub::DeviceFor::ForEachInLayout(d_temp_storage,
+                                                                      temp_storage_bytes,
+                                                                      layout_mapping,
+                                                                      op,
+                                                                      stream));
+    }
 };
 
 END_HIPCUB_NAMESPACE