Develop stream 2025-01-23

.gitlab-ci.yml

@@ -1,6 +1,6 @@
 # MIT License
 #
-# Copyright (c) 2017-2024 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (c) 2017-2025 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
@@ -42,6 +42,17 @@ stages:
   - test
   - benchmark
 
+workflow:
+  rules:
+    - if: $CI_MERGE_REQUEST_LABELS =~ /CI Skip/
+      when: never
+    - if: $CI_MERGE_REQUEST_TITLE !~ /Draft:/
+      variables:
+        ROCPRIM_TEST_RUNS: 3
+    - if: $CI_MERGE_REQUEST_TITLE =~ /Draft:/
+      variables:
+        ROCPRIM_TEST_RUNS: 1
+
 variables:
   PACKAGE_DIR: $BUILD_DIR/package
   AUTOTUNE_CONFIG_DIR: ${CI_PROJECT_DIR}/autotune_config
@@ -110,8 +121,7 @@ copyright-date:
       -D BUILD_TEST=ON
       -D BUILD_EXAMPLE=ON
       -D BUILD_BENCHMARK=ON
-      -D GPU_TARGETS=$GPU_TARGETS
-      -D AMDGPU_TEST_TARGETS=$GPU_TARGETS
+      -D AMDGPU_TARGETS=$GPU_TARGETS
       -D CMAKE_C_COMPILER_LAUNCHER=phc_sccache_c
       -D CMAKE_CXX_COMPILER_LAUNCHER=phc_sccache_cxx
       -D CMAKE_CXX_STANDARD=14
@@ -170,7 +180,7 @@ build:cmake-minimum-apt:
     - | # Add hardened libc++ assertions for tests only
         if [[ $BUILD_TARGET == "TEST" ]]; then
             echo "Configuring with hardened libc++!"
-            EXTRA_CMAKE_CXX_FLAGS+=" -D_GLIBCXX_ASSERTIONS=ON"
+            EXTRA_CMAKE_CXX_FLAGS+=" -D_GLIBCXX_ASSERTIONS=ON -D ROCPRIM_ENABLE_ASSERTS=ON"
         fi
     - cmake
       -G Ninja
@@ -181,8 +191,7 @@ build:cmake-minimum-apt:
       -D CMAKE_C_COMPILER_LAUNCHER=phc_sccache_c
       -D CMAKE_CXX_COMPILER_LAUNCHER=phc_sccache_cxx
       -D BUILD_EXAMPLE=ON
-      -D GPU_TARGETS=$GPU_TARGETS
-      -D AMDGPU_TEST_TARGETS=$GPU_TARGETS
+      -D AMDGPU_TARGETS=$GPU_TARGETS
       -D CMAKE_CXX_STANDARD="$BUILD_VERSION"
       -S $CI_PROJECT_DIR
       -B $BUILD_DIR
@@ -198,8 +207,9 @@ build:cmake-minimum-apt:
       - $BUILD_DIR/test/CTestTestfile.cmake
       - $BUILD_DIR/test/rocprim/CTestTestfile.cmake
       - $BUILD_DIR/test/rocprim/test_*
+      - $BUILD_DIR/test/rocprim/libtest_*
       - $BUILD_DIR/test/test_*
-    expire_in: 2 weeks
+    expire_in: 1 day
 
 build:cmake-latest:
   stage: build
@@ -249,7 +259,7 @@ build:package:
     paths:
       - $PACKAGE_DIR/rocprim*.deb
       - $PACKAGE_DIR/rocprim*.zip
-    expire_in: 2 weeks
+    expire_in: 1 day
 
 build:windows:
   stage: build
@@ -269,7 +279,7 @@ build:windows:
       -S $CI_PROJECT_DIR
       -B $CI_PROJECT_DIR/build
       -D BUILD_$BUILD_TARGET=ON
-      -D GPU_TARGETS=$GPU_TARGET
+      -D AMDGPU_TARGETS=$GPU_TARGET
       -D CMAKE_CXX_COMPILER:PATH="${env:HIP_PATH}\bin\clang++.exe"
       -D CMAKE_PREFIX_PATH:PATH="${env:HIP_PATH}"
       -D CMAKE_BUILD_TYPE="$BUILD_TYPE"
@@ -285,7 +295,7 @@ build:windows:
       - $CI_PROJECT_DIR/build/CMakeCache.txt
       - $CI_PROJECT_DIR/build/.ninja_log
       - $CI_PROJECT_DIR/build/CTestTestfile.cmake
-    expire_in: 2 weeks
+    expire_in: 1 day
 
 autotune:build:
   stage: autotune
@@ -317,7 +327,7 @@ autotune:build:
       -D BUILD_EXAMPLE=OFF
       -D BUILD_BENCHMARK=ON
       -D BENCHMARK_CONFIG_TUNING=ON
-      -D GPU_TARGETS=$GPU_TARGETS
+      -D AMDGPU_TARGETS=$GPU_TARGETS
       -D CMAKE_C_COMPILER_LAUNCHER=phc_sccache_c
       -D CMAKE_CXX_COMPILER_LAUNCHER=phc_sccache_cxx
       -D CMAKE_CXX_STANDARD=14
@@ -334,12 +344,13 @@ autotune:build:
       - $BUILD_DIR/deps/googlebenchmark/
     expire_in: 1 week
 
-test:
+.test:common:
   stage: test
+  tags:
+    - rocm
+    - $GPU
   extends:
     - .cmake-minimum
-    - .rules:test
-    - .gpus:rocm
   needs:
     - job: build:cmake-minimum
       parallel:
@@ -359,10 +370,32 @@ test:
     - HSA_ENABLE_SDMA=0 ctest
       --output-on-failure
       --repeat-until-fail 2
-      --tests-regex "hip|$GPU_TARGET"
       --resource-spec-file ./resources.json
       --parallel $PARALLEL_JOBS
 
+test:any-gpu:
+  variables:
+    GPU: ""
+    PARALLEL_JOBS: 1
+  extends:
+    - .test:common
+  rules:
+    - if: $CI_MERGE_REQUEST_TITLE =~ /Draft:/ && $CI_MERGE_REQUEST_LABELS !~ /Arch::/
+
+test:label-arch:
+  extends:
+    - .gpus:rocm
+    - .test:common
+    - .rules:arch-labels
+
+test:all-gpus:
+  variables:
+    SHOULD_BE_UNDRAFTED: "true"
+  extends:
+    - .gpus:rocm
+    - .test:common
+    - .rules:test
+
 .test-windows-base:
   stage: test
   extends:
@@ -402,7 +435,7 @@ test-windows-release:
       -G Ninja
       -D CMAKE_CXX_COMPILER="$AMDCLANG"
       -D CMAKE_BUILD_TYPE=Release
-      -D GPU_TARGETS=$GPU_TARGETS
+      -D AMDGPU_TARGETS=$GPU_TARGETS
       -D CMAKE_CXX_STANDARD=14
       -S "$CI_PROJECT_DIR/test/extra"
       -B "$CI_PROJECT_DIR/package_test"
@@ -428,7 +461,8 @@ test:install:
       -D CMAKE_CXX_STANDARD=14
       -B build
       -S $CI_PROJECT_DIR
-    - $SUDO_CMD cmake --build build --target install
+    # Preserve $PATH when sudoing
+    - $SUDO_CMD env PATH="$PATH" cmake --build build --target install
     - !reference [.test-package, script]
 
 test:deb:
@@ -452,6 +486,10 @@ test:docs:
   extends:
     - .rules:test
     - .build:docs
+  artifacts:
+    paths:
+      - $DOCS_DIR/_build/html/
+    expire_in: 2 weeks
 
 .benchmark-base:
   stage: benchmark

CHANGELOG.md

@@ -5,7 +5,30 @@ Full documentation for rocPRIM is available at [https://rocm.docs.amd.com/projec
 ## rocPRIM 3.5.0 for ROCm 6.5.0
 
 ### Added
-* gfx950 support
+
+* Added gfx950 support.
+* Added `rocprim::key_value_pair::operator==`.
+* Added the `rocprim::unrolled_copy` thread function to copy multiple items inside a thread.
+* Added the `rocprim::unrolled_thread_load` function to load multiple items inside a thread using `rocprim::thread_load`.
+* Added `rocprim::int128_t` and `rocprim::uint128_t` to benchmarks for improved performance evaluation on 128-bit integers.
+* Added `rocprim::int128_t` to the supported autotuning types to improve performance for 128-bit integers.
+* Added the `rocprim::merge_inplace` function for merging in-place.
+* Added initial value support for warp- and block-level inclusive scan.
+
+### Changed
+
+* A new version of `rocprim::thread_load` and `rocprim::thread_store` replace the deprecated `rocprim::thread_load` and `rocprim::thread_store` functions. The versions avoid inline assembly where possible, and don't hinder the optimizer as much as a result.
+* Renamed `rocprim::load_cs` to `rocprim::load_nontemporal` and `rocprim::store_cs` to `rocprim::store_nontemporal` to express the intent of these load and store methods better.
+* All kernels now have hidden symbol visibility. All symbols now have inline namespaces that include the library version, for example, `rocprim::ROCPRIM_300400_NS::symbol` instead of `rocPRIM::symbol`, letting the user link multiple libraries built with different versions of rocPRIM.
+
+### Resolved issues
+
+* Fixed device radix sort not returning the correct required temporary storage when a double buffer contains `nullptr`.
+* Fixed constness of equality operators (`==` and `!=`) in `rocprim::key_value_pair`.
+
+### Deprecations
+
+* `rocprim::load_cs` and `rocprim::store_cs` are deprecated. Use `rocprim::load_nontemporal` and `rocprim::store_nontemporal` now.
 
 ## rocPRIM 3.4.0 for ROCm 6.4.0
 
@@ -26,6 +49,10 @@ Full documentation for rocPRIM is available at [https://rocm.docs.amd.com/projec
 * Added the parallel `search` and `find_end` device functions similar to `std::search` and `std::find_end`, these functions search for the first and last occurrence of the sequence respectively.
 * Added a parallel device-level function, `rocprim::search_n`, similar to the C++ Standard Library `std::search_n` algorithm.
 * Added new constructors and a `base` function, and added `constexpr` specifier to all functions in `rocprim::reverse_iterator` to improve parity with the C++17 `std::reverse_iterator`.
+* Added hipGraph support to device run-length-encode for nontrivial runs (`rocprim::run_length_encode_non_trivial_runs`).
+* Added configuration autotuning to device run-length-encode for nontrivial runs (`rocprim::run_length_encode_non_trivial_runs`) for improved performance on selected architectures.
+* Added configuration autotuning to device run-length-encode for trivial runs (`rocprim::run_length_encode`) for improved performance on selected architectures.
+* Added a new type traits interface to enable users to provide additional type trait information to rocPRIM, facilitating better compatibility with custom types.
 
 ### Changed
 
@@ -36,6 +63,7 @@ Full documentation for rocPRIM is available at [https://rocm.docs.amd.com/projec
 
 * Removed HIP-CPU support. HIP-CPU support was experimental and broken.
 * Changed the C++ version from 14 to 17. C++14 will be deprecated in the next major release.
+* You can use CMake HIP language support with CMake 3.18 and later. To use HIP language support, run `cmake` with `-DUSE_HIPCXX=ON` instead of setting the `CXX` variable to the path to a HIP-aware compiler.
 
 ### Resolved issues
 

CMakeLists.txt

@@ -1,6 +1,6 @@
 # MIT License
 #
-# Copyright (c) 2017-2024 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (c) 2017-2025 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
@@ -74,6 +74,7 @@ option(BUILD_NAIVE_BENCHMARK "Build naive benchmarks" OFF)
 cmake_dependent_option(BUILD_DOCS "Build documentation (requires sphinx)" OFF "NOT ONLY_INSTALL" OFF)
 option(BUILD_CODE_COVERAGE "Build with code coverage enabled" OFF)
 option(ROCPRIM_INSTALL "Enable installation of rocPRIM (projects embedding rocPRIM may want to turn this OFF)" ON)
+option(ROCPRIM_ENABLE_ASSERTS "Enable asserts in release build)" OFF)
 
 check_language(HIP)
 cmake_dependent_option(USE_HIPCXX "Use CMake HIP language support" OFF CMAKE_HIP_COMPILER OFF)
@@ -104,6 +105,21 @@ if(NOT CMAKE_BUILD_TYPE AND NOT CMAKE_CONFIGURATION_TYPES)
   set_property(CACHE CMAKE_BUILD_TYPE PROPERTY STRINGS "" "Debug" "Release" "MinSizeRel" "RelWithDebInfo")
 endif()
 
+if(ROCPRIM_ENABLE_ASSERTS)
+  if(NOT "${CMAKE_BUILD_TYPE}" STREQUAL "")
+      string(TOUPPER ${CMAKE_BUILD_TYPE} BUILD_TYPE)
+      set(BUILD_TYPE_CXX_FLAGS "CMAKE_CXX_FLAGS_${BUILD_TYPE}")
+      set(BUILD_TYPE_C_FLAGS "CMAKE_C_FLAGS_${BUILD_TYPE}")
+  endif()
+
+  string(REGEX REPLACE "-DNDEBUG( |$)" "" CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS}")
+  string(REGEX REPLACE "-DNDEBUG( |$)" "" CMAKE_C_FLAGS "${CMAKE_C_FLAGS}")
+  if(NOT "${CMAKE_BUILD_TYPE}" STREQUAL "")
+      string(REGEX REPLACE "-DNDEBUG( |$)" "" ${BUILD_TYPE_CXX_FLAGS} "${${BUILD_TYPE_CXX_FLAGS}}")
+      string(REGEX REPLACE "-DNDEBUG( |$)" "" ${BUILD_TYPE_C_FLAGS} "${${BUILD_TYPE_C_FLAGS}}")
+  endif()
+endif()
+
 set(CMAKE_INSTALL_RPATH_USE_LINK_PATH TRUE CACHE BOOL "Add paths to linker search and installed rpath")
 
 if(DEFINED BUILD_SHARED_LIBS)
@@ -167,6 +183,7 @@ endif()
 # Setup VERSION
 set(VERSION_STRING "3.4.0")
 rocm_setup_version(VERSION ${VERSION_STRING})
+math(EXPR rocprim_VERSION_NUMBER "${rocprim_VERSION_MAJOR} * 100000 + ${rocprim_VERSION_MINOR} * 100 + ${rocprim_VERSION_PATCH}")
 
 # Print configuration summary
 include(cmake/Summary.cmake)