ROCm · Nemi26 · Oct 9, 2025 · Oct 9, 2025 · Oct 9, 2025 · Oct 9, 2025
@@ -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'
+]
+
 ---
@@ -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.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.
+* Benchmarking now uses primbench for its benchmarks instead of Google Benchmark.
+  * See `shared/primbench/README.md` for its documentation.
+
+### 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

@@ -1,6 +1,6 @@
 # MIT License
 #
-# Copyright (c) 2017-2025 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (c) 2017-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,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

@@ -30,9 +30,9 @@ environment, hipCUB uses the rocPRIM library as the backend.
 Optional:
 
 * [GoogleTest](https://github.com/google/googletest)
-* [Google Benchmark](https://github.com/google/benchmark)
-
-GoogleTest and Google Benchmark are automatically downloaded and built by the CMake script.
+  * This is automatically downloaded and built by the CMake script
+* [AMD SMI](https://github.com/ROCm/amdsmi)
+  * Required only for benchmarks. Building benchmarks is off by default.
 
 ## Build and install
 
@@ -69,7 +69,7 @@ mkdir build; cd build
 #   BUILD_TEST                   - OFF by default,
 #   BUILD_BENCHMARK              - OFF by default.
 #   ROCPRIM_FETCH_METHOD         - One of PACKAGE (default), DOWNLOAD, and MONOREPO. See below for a description of each.
-#   EXTERNAL_DEPS_FORCE_DOWNLOAD - OFF by default, forces download for non-ROCm dependencies (eg. Google Test / Benchmark).
+#   EXTERNAL_DEPS_FORCE_DOWNLOAD - OFF by default, forces download for non-ROCm dependencies (e.g., Google Test).
 #   DOWNLOAD_CUB                 - OFF by default, (Nvidia CUB backend only) forces download of CUB instead of searching for an installed package.
 #   BUILD_OFFLOAD_COMPRESS       - ON by default, compresses device code to reduce the size of the generated binary.
 #   BUILD_EXAMPLE                - OFF by default, builds examples.
@@ -253,24 +253,20 @@ ctest --resource-spec-file <path-to-your-resources.json> --parallel <number-of-j
 
 ## Running benchmarks
 
+hipCUB uses [primbench](https://github.com/ROCm/rocm-libraries/tree/develop/shared/primbench) for benchmarking; see its [command-line options](https://github.com/ROCm/rocm-libraries/tree/develop/shared/primbench#command-line-options) for advanced usage.
+
 ```shell
 # Go to hipCUB build directory
 cd projects/hipcub; cd build
 
 # To run benchmark for warp functions:
-# Further option can be found using --help
-# [] Fields are optional
-./benchmark/benchmark_warp_<function_name> [--size <size>] [--trials <trials>]
+./benchmark/benchmark_warp_<function_name>
 
 # To run benchmark for block functions:
-# Further option can be found using --help
-# [] Fields are optional
-./benchmark/benchmark_block_<function_name> [--size <size>] [--trials <trials>]
+./benchmark/benchmark_block_<function_name>
 
 # To run benchmark for device functions:
-# Further option can be found using --help
-# [] Fields are optional
-./benchmark/benchmark_device_<function_name> [--size <size>] [--trials <trials>]
+./benchmark/benchmark_device_<function_name>
 ```
 
 ## Building the documentation locally

@@ -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
@@ -20,52 +20,79 @@
 # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
 # SOFTWARE.
 
+# Force CMake to find Git immediately so ${GIT_EXECUTABLE} is populated
+find_package(Git REQUIRED)
+
+# Get the current branch, or fall back to DETACHED
+execute_process(
+  COMMAND ${GIT_EXECUTABLE} symbolic-ref -q --short HEAD
+  WORKING_DIRECTORY ${CMAKE_CURRENT_LIST_DIR}
+  OUTPUT_VARIABLE BRANCH_NAME
+  OUTPUT_STRIP_TRAILING_WHITESPACE
+)
+if(BRANCH_NAME STREQUAL "")
+  set(BRANCH_NAME "DETACHED")
+endif()
+
+# Get the short commit hash
+execute_process(
+  COMMAND ${GIT_EXECUTABLE} rev-parse --short HEAD
+  WORKING_DIRECTORY ${CMAKE_CURRENT_LIST_DIR}
+  OUTPUT_VARIABLE COMMIT_HASH
+  OUTPUT_STRIP_TRAILING_WHITESPACE
+)
+
 function(add_hipcub_benchmark BENCHMARK_SOURCE)
   get_filename_component(BENCHMARK_TARGET ${BENCHMARK_SOURCE} NAME_WE)
+
   if(USE_HIPCXX)
     set_source_files_properties(${BENCHMARK_SOURCE} PROPERTIES LANGUAGE HIP)
   endif()
+
   add_executable(${BENCHMARK_TARGET} ${BENCHMARK_SOURCE})
-  target_include_directories(${BENCHMARK_TARGET} SYSTEM BEFORE
-    PUBLIC
-      "$<BUILD_INTERFACE:${GOOGLEBENCHMARK_ROOT}/include>"
-  )
-  target_link_libraries(${BENCHMARK_TARGET}
-    PRIVATE
-      benchmark::benchmark
-      hipcub
-  )
-  if((HIP_COMPILER STREQUAL "nvcc"))
+  target_link_libraries(${BENCHMARK_TARGET} PRIVATE hipcub)
+
+  target_compile_definitions(${BENCHMARK_TARGET} PUBLIC BRANCH_NAME="${BRANCH_NAME}")
+  target_compile_definitions(${BENCHMARK_TARGET} PUBLIC COMMIT_HASH="${COMMIT_HASH}")
+
+  target_include_directories(${BENCHMARK_TARGET} PRIVATE ../../../shared/primbench)
+
+  if(WIN32)
+    target_compile_definitions(${BENCHMARK_TARGET} PRIVATE PRIMBENCH_NO_MONITORING)
+  elseif((HIP_COMPILER STREQUAL "nvcc"))
     set_property(TARGET ${BENCHMARK_TARGET} PROPERTY CUDA_STANDARD 17)
     set_source_files_properties(${BENCHMARK_SOURCE} PROPERTIES LANGUAGE CUDA)
     target_compile_options(${BENCHMARK_TARGET}
       PRIVATE
-        $<$<COMPILE_LANGUAGE:CUDA>:--expt-extended-lambda>
-    )
-    target_link_libraries(${BENCHMARK_TARGET}
-      PRIVATE
-        hipcub_cub
+      $<$<COMPILE_LANGUAGE:CUDA>:--expt-extended-lambda>
     )
+    target_link_libraries(${BENCHMARK_TARGET} PRIVATE hipcub_cub nvidia-ml)
+  else()
+    target_link_libraries(${BENCHMARK_TARGET} PRIVATE amd_smi)
   endif()
+
   set_target_properties(${BENCHMARK_TARGET}
     PROPERTIES
-      RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/benchmark"
+    RUNTIME_OUTPUT_DIRECTORY "${CMAKE_BINARY_DIR}/benchmark"
   )
   rocm_install(TARGETS ${BENCHMARK_TARGET} COMPONENT benchmarks)
-  if (WIN32 AND NOT DEFINED DLLS_COPIED)
+
+  if(WIN32 AND NOT DEFINED DLLS_COPIED)
     set(DLLS_COPIED "YES")
     set(DLLS_COPIED ${DLLS_COPIED} PARENT_SCOPE)
+
     # for now adding in all .dll as dependency chain is not cmake based on win32
-    file( GLOB third_party_dlls
-    LIST_DIRECTORIES ON
-    CONFIGURE_DEPENDS
-    ${HIP_DIR}/bin/*.dll
-    ${CMAKE_SOURCE_DIR}/rtest.*
+    file(GLOB third_party_dlls
+      LIST_DIRECTORIES ON
+      CONFIGURE_DEPENDS
+      ${HIP_DIR}/bin/*.dll
+      ${CMAKE_SOURCE_DIR}/rtest.*
     )
-    foreach( file_i ${third_party_dlls})
-      add_custom_command( TARGET ${BENCHMARK_TARGET} POST_BUILD COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${file_i} ${PROJECT_BINARY_DIR}/benchmark )
-    endforeach( file_i )
-  endif()         
+
+    foreach(file_i ${third_party_dlls})
+      add_custom_command(TARGET ${BENCHMARK_TARGET} POST_BUILD COMMAND ${CMAKE_COMMAND} ARGS -E copy_if_different ${file_i} ${PROJECT_BINARY_DIR}/benchmark)
+    endforeach(file_i)
+  endif()
 endfunction()
 
 # ****************************************************************************
@@ -76,34 +103,33 @@ add_hipcub_benchmark(benchmark_block_discontinuity.cpp)
 add_hipcub_benchmark(benchmark_block_exchange.cpp)
 add_hipcub_benchmark(benchmark_block_histogram.cpp)
 add_hipcub_benchmark(benchmark_block_merge_sort.cpp)
-add_hipcub_benchmark(benchmark_block_radix_sort.cpp)
 add_hipcub_benchmark(benchmark_block_radix_rank.cpp)
+add_hipcub_benchmark(benchmark_block_radix_sort.cpp)
 add_hipcub_benchmark(benchmark_block_reduce.cpp)
 add_hipcub_benchmark(benchmark_block_run_length_decode.cpp)
 add_hipcub_benchmark(benchmark_block_scan.cpp)
 add_hipcub_benchmark(benchmark_block_shuffle.cpp)
 add_hipcub_benchmark(benchmark_device_adjacent_difference.cpp)
 add_hipcub_benchmark(benchmark_device_batch_copy.cpp)
 add_hipcub_benchmark(benchmark_device_batch_memcpy.cpp)
-add_hipcub_benchmark(benchmark_device_for.cpp)
+add_hipcub_benchmark(benchmark_device_for_each.cpp)
 add_hipcub_benchmark(benchmark_device_histogram.cpp)
 add_hipcub_benchmark(benchmark_device_memory.cpp)
-add_hipcub_benchmark(benchmark_device_merge.cpp)
 add_hipcub_benchmark(benchmark_device_merge_sort.cpp)
+add_hipcub_benchmark(benchmark_device_merge.cpp)
 add_hipcub_benchmark(benchmark_device_partition.cpp)
 add_hipcub_benchmark(benchmark_device_radix_sort.cpp)
 add_hipcub_benchmark(benchmark_device_reduce_by_key.cpp)
 add_hipcub_benchmark(benchmark_device_reduce.cpp)
 add_hipcub_benchmark(benchmark_device_run_length_encode.cpp)
 add_hipcub_benchmark(benchmark_device_scan.cpp)
-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_segmented_sort.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_merge_sort.cpp)
 add_hipcub_benchmark(benchmark_warp_reduce.cpp)
 add_hipcub_benchmark(benchmark_warp_scan.cpp)
 add_hipcub_benchmark(benchmark_warp_store.cpp)
-add_hipcub_benchmark(benchmark_warp_merge_sort.cpp)