Develop stream 2024-10-29

.gitlab-ci.yml

@@ -114,6 +114,7 @@ copyright-date:
       -D AMDGPU_TEST_TARGETS=$GPU_TARGETS
       -D CMAKE_C_COMPILER_LAUNCHER=phc_sccache_c
       -D CMAKE_CXX_COMPILER_LAUNCHER=phc_sccache_cxx
+      -D CMAKE_CXX_STANDARD=14
       -S $CI_PROJECT_DIR
       -B $BUILD_DIR
     - cmake
@@ -182,6 +183,7 @@ build:cmake-minimum-apt:
       -D BUILD_EXAMPLE=ON
       -D GPU_TARGETS=$GPU_TARGETS
       -D AMDGPU_TEST_TARGETS=$GPU_TARGETS
+      -D CMAKE_CXX_STANDARD="$BUILD_VERSION"
       -S $CI_PROJECT_DIR
       -B $BUILD_DIR
     - cmake --build $BUILD_DIR
@@ -210,6 +212,7 @@ build:cmake-latest:
     matrix:
       - BUILD_TYPE: Release
         BUILD_TARGET: [BENCHMARK, TEST]
+        BUILD_VERSION: [14, 17]
 
 build:cmake-minimum:
   needs: []
@@ -220,6 +223,7 @@ build:cmake-minimum:
     matrix:
       - BUILD_TYPE: [Debug, Release]
         BUILD_TARGET: [BENCHMARK, TEST]
+        BUILD_VERSION: 14
 
 build:package:
   stage: build
@@ -236,6 +240,7 @@ build:package:
       -G Ninja
       -D CMAKE_CXX_COMPILER="$AMDCLANG"
       -D CMAKE_BUILD_TYPE=Release
+      -D CMAKE_CXX_STANDARD=14
       -B $PACKAGE_DIR
       -S $CI_PROJECT_DIR
     - cd $PACKAGE_DIR
@@ -268,6 +273,7 @@ build:windows:
       -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"
+      -D CMAKE_CXX_STANDARD=14
     - cmake --build "$CI_PROJECT_DIR/build"
   artifacts:
     paths:
@@ -314,6 +320,7 @@ autotune:build:
       -D GPU_TARGETS=$GPU_TARGETS
       -D CMAKE_C_COMPILER_LAUNCHER=phc_sccache_c
       -D CMAKE_CXX_COMPILER_LAUNCHER=phc_sccache_cxx
+      -D CMAKE_CXX_STANDARD=14
     - cmake --build . --target $BENCHMARK_TARGETS
     - 'rm -rf $BUILD_DIR/benchmark/benchmark*.parallel'
     # The autotune benchmarks get very large, above GitLabs upload limit. Fortunately they compress well.
@@ -339,6 +346,7 @@ test:
         matrix:
           - BUILD_TYPE: Release
             BUILD_TARGET: TEST
+            BUILD_VERSION: 14
   script:
     - cd $BUILD_DIR
     - cmake
@@ -395,6 +403,7 @@ test-windows-release:
       -D CMAKE_CXX_COMPILER="$AMDCLANG"
       -D CMAKE_BUILD_TYPE=Release
       -D GPU_TARGETS=$GPU_TARGETS
+      -D CMAKE_CXX_STANDARD=14
       -S "$CI_PROJECT_DIR/test/extra"
       -B "$CI_PROJECT_DIR/package_test"
     - cmake --build "$CI_PROJECT_DIR/package_test"
@@ -416,6 +425,7 @@ test:install:
       -G Ninja
       -D CMAKE_CXX_COMPILER="$AMDCLANG"
       -D CMAKE_BUILD_TYPE=Release
+      -D CMAKE_CXX_STANDARD=14
       -B build
       -S $CI_PROJECT_DIR
     - $SUDO_CMD cmake --build build --target install
@@ -458,6 +468,7 @@ benchmark:
         matrix:
           - BUILD_TYPE: Release
             BUILD_TARGET: BENCHMARK
+            BUILD_VERSION: 14
   extends:
     - .cmake-minimum
     - .gpus:rocm

CHANGELOG.md

@@ -11,18 +11,29 @@ Full documentation for rocPRIM is available at [https://rocm.docs.amd.com/projec
 * Added the parallel `find_first_of` device function with autotuned configurations, this function is similar to `std::find_first_of`, it searches for the first occurrence of any of the provided elements.
 * Added `--emulation` option added for `rtest.py`
   * Unit tests can be run with `[--emulation|-e|--test|-t]=<test_name>`
+* Added a parallel device-level function, `rocprim::adjacent_find`, similar to the C++ Standard Library `std::adjacent_find` algorithm.
+* Added configuration autotuning to device adjacent find (`rocprim::adjacent_find`) for improved performance on selected architectures.
+* Added rocprim::numeric_limits which is an extension of `std::numeric_limits`, which includes support for 128-bit integers.
+* Added rocprim::int128_t and rocprim::uint128_t which are the __int128_t and __uint128_t types.
+* 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`.
 
 ### Changed
 * Changed the subset of tests that are run for smoke tests such that the smoke test will complete with faster run-time and to never exceed 2GB of vram usage. Use `python rtest.py [--emulation|-e|--test|-t]=smoke` to run these tests.
 * The `rtest.py` options have changed. `rtest.py` is now run with at least either `--test|-t` or `--emulation|-e`, but not both options.
 
+* 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.
+
 ### Resolved issues
 * Fixed an issue where `rmake.py` would generate wrong CMAKE commands while using Linux enviorment
 * Fixed an issue where `rocprim::partial_sort_copy` would yield a compile error if the input iterator is const.
 * Fixed incorrect 128-bit signed and unsigned integers type traits.
 * Fixed compilation issue when `rocprim::radix_key_codec<...>` is specialized with a 128-bit integer.
 
 ### Upcoming changes
+* Using the initialisation constructor of `rocprim::reverse_iterator` will throw a deprecation warning. It will be marked as explicit in the next major release.
 
 ## rocPRIM 3.3.0 for ROCm 6.3.0
 

CMakeLists.txt

@@ -48,7 +48,6 @@ option(BUILD_BENCHMARK "Build benchmarks" OFF)
 option(BUILD_NAIVE_BENCHMARK "Build naive benchmarks" OFF)
 option(BUILD_EXAMPLE "Build examples" OFF)
 option(BUILD_DOCS "Build documentation (requires sphinx)" OFF)
-option(USE_HIP_CPU "Prefer HIP-CPU runtime instead of HW acceleration" OFF)
 # Disables building tests, benchmarks, examples
 option(ONLY_INSTALL "Only install" OFF)
 option(BUILD_CODE_COVERAGE "Build with code coverage enabled" OFF)
@@ -70,50 +69,57 @@ endif()
 set(CMAKE_INSTALL_RPATH_USE_LINK_PATH TRUE CACHE BOOL "Add paths to linker search and installed rpath")
 
 # Set CXX flags
-set(CMAKE_CXX_STANDARD 14)
+if (NOT DEFINED CMAKE_CXX_STANDARD)
+  set(CMAKE_CXX_STANDARD 17)
+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")
+endif()
+
 if(DEFINED BUILD_SHARED_LIBS)
   set(PKG_BUILD_SHARED_LIBS ${BUILD_SHARED_LIBS})
-else()
+else()  
   set(PKG_BUILD_SHARED_LIBS ON)
-endif()
+endif()  
 set(BUILD_SHARED_LIBS OFF) # don't build client dependencies as shared
-if(NOT USE_HIP_CPU)
-  # Get dependencies (required here to get rocm-cmake)
-  include(cmake/Dependencies.cmake)
-  # Use target ID syntax if supported for GPU_TARGETS
-  if (NOT DEFINED AMDGPU_TARGETS)
-    set(GPU_TARGETS "all" CACHE STRING "GPU architectures to compile for")
+
+# Get dependencies (required here to get rocm-cmake)
+include(cmake/Dependencies.cmake)
+# Use target ID syntax if supported for GPU_TARGETS
+if (NOT DEFINED AMDGPU_TARGETS)
+  set(GPU_TARGETS "all" CACHE STRING "GPU architectures to compile for")
+else()
+  set(GPU_TARGETS "${AMDGPU_TARGETS}" CACHE STRING "GPU architectures to compile for")
+endif()
+set_property(CACHE GPU_TARGETS PROPERTY STRINGS "all")
+
+if(GPU_TARGETS STREQUAL "all")
+  if(BUILD_ADDRESS_SANITIZER)
+    # ASAN builds require xnack
+    rocm_check_target_ids(DEFAULT_AMDGPU_TARGETS
+      TARGETS "gfx908:xnack+;gfx90a:xnack+;gfx940:xnack+;gfx941:xnack+;gfx942:xnack+"
+    )
   else()
-    set(GPU_TARGETS "${AMDGPU_TARGETS}" CACHE STRING "GPU architectures to compile for")
-  endif()
-  set_property(CACHE GPU_TARGETS PROPERTY STRINGS "all")
-
-  if(GPU_TARGETS STREQUAL "all")
-    if(BUILD_ADDRESS_SANITIZER)
-      # ASAN builds require xnack
-      rocm_check_target_ids(DEFAULT_AMDGPU_TARGETS
-        TARGETS "gfx908:xnack+;gfx90a:xnack+;gfx940:xnack+;gfx941:xnack+;gfx942:xnack+"
-      )
-    else()
-      rocm_check_target_ids(DEFAULT_AMDGPU_TARGETS
-        TARGETS "gfx803;gfx900:xnack-;gfx906:xnack-;gfx908:xnack-;gfx90a:xnack-;gfx90a:xnack+;gfx940;gfx941;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1151;gfx1200;gfx1201"
-      )
-    endif()
-
-    set(GPU_TARGETS "${DEFAULT_AMDGPU_TARGETS}" CACHE STRING "GPU architectures to compile for" FORCE)
+    rocm_check_target_ids(DEFAULT_AMDGPU_TARGETS
+      TARGETS "gfx803;gfx900:xnack-;gfx906:xnack-;gfx908:xnack-;gfx90a:xnack-;gfx90a:xnack+;gfx940;gfx941;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1151;gfx1200;gfx1201"
+    )
   endif()
 
-  # TODO: Fix VerifyCompiler for HIP on Windows
-  if (NOT WIN32)
-    include(cmake/VerifyCompiler.cmake)
-  endif()
-  list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH} ${ROCM_PATH}/hip ${ROCM_PATH}/llvm ${ROCM_ROOT}/llvm ${ROCM_ROOT} ${ROCM_ROOT}/hip)
-  find_package(hip REQUIRED CONFIG PATHS ${HIP_DIR} ${ROCM_PATH} /opt/rocm)
+  set(GPU_TARGETS "${DEFAULT_AMDGPU_TARGETS}" CACHE STRING "GPU architectures to compile for" FORCE)
 endif()
 
+# TODO: Fix VerifyCompiler for HIP on Windows
+if (NOT WIN32)
+  include(cmake/VerifyCompiler.cmake)
+endif()
+list(APPEND CMAKE_PREFIX_PATH ${ROCM_PATH} ${ROCM_PATH}/hip ${ROCM_PATH}/llvm ${ROCM_ROOT}/llvm ${ROCM_ROOT} ${ROCM_ROOT}/hip)
+find_package(hip REQUIRED CONFIG PATHS ${HIP_DIR} ${ROCM_PATH} /opt/rocm)
+
 # FOR HANDLING ENABLE/DISABLE OPTIONAL BACKWARD COMPATIBILITY for FILE/FOLDER REORG
 option(BUILD_FILE_REORG_BACKWARD_COMPATIBILITY "Build with file/folder reorg with backward compatibility enabled" OFF)
 if(ROCPRIM_INSTALL AND BUILD_FILE_REORG_BACKWARD_COMPATIBILITY AND NOT WIN32)
@@ -130,11 +136,6 @@ if(BUILD_CODE_COVERAGE)
   add_link_options(--coverage)
 endif()
 
-if(USE_HIP_CPU)
-  # Get dependencies
-  include(cmake/Dependencies.cmake)
-endif()
-
 # Setup VERSION
 set(VERSION_STRING "3.3.0")
 rocm_setup_version(VERSION ${VERSION_STRING})

README.md

@@ -14,7 +14,7 @@ develop performant GPU-accelerated code on AMD ROCm platforms.
   * Including
     [HIP-clang](https://github.com/ROCm/HIP/blob/master/INSTALL.md#hip-clang)
     compiler
-* C++14
+* C++17
 * Python 3.6 or higher (HIP on Windows only, required only for install script)
 * Visual Studio 2019 with Clang support (HIP on Windows only)
 * Strawberry Perl (HIP on Windows only)
@@ -110,11 +110,6 @@ You can build and install rocPRIM on Linux or Windows.
   # before 'cmake' or setting cmake option 'CMAKE_CXX_COMPILER' to path to the compiler.
   # Using HIP-clang:
   [CXX=hipcc] cmake -DBUILD_BENCHMARK=ON ../.
-  #
-  # ! EXPERIMENTAL !
-  # Alternatively one may build using the experimental (and highly incomplete) HIP-CPU back-end for host-side
-  # execution using any C++17 conforming compiler (supported by HIP-CPU). AMDGPU_* options are unavailable in this case. 
-  #   USE_HIP_CPU - OFF by default
 
   # Build
   make -j4

benchmark/CMakeLists.txt

@@ -68,24 +68,10 @@ function(add_rocprim_benchmark BENCHMARK_SOURCE)
       rocprim
       benchmark::benchmark
   )
-  if(NOT USE_HIP_CPU)
-    target_link_libraries(${BENCHMARK_TARGET}
-      PRIVATE
-        rocprim_hip
-    )
-  else()
-    target_link_libraries(${BENCHMARK_TARGET}
-      PRIVATE
-        Threads::Threads
-        hip_cpu_rt::hip_cpu_rt
-    )
-    if(STL_DEPENDS_ON_TBB)
-      target_link_libraries(${BENCHMARK_TARGET}
-        PRIVATE
-          TBB::tbb
-      )
-    endif()
-  endif()
+  target_link_libraries(${BENCHMARK_TARGET}
+    PRIVATE
+      rocprim_hip
+  )
 
   target_compile_options(${BENCHMARK_TARGET}
     PRIVATE
@@ -134,9 +120,11 @@ add_rocprim_benchmark(benchmark_block_scan.cpp)
 add_rocprim_benchmark(benchmark_block_sort.cpp)
 add_rocprim_benchmark(benchmark_config_dispatch.cpp)
 add_rocprim_benchmark(benchmark_device_adjacent_difference.cpp)
+add_rocprim_benchmark(benchmark_device_adjacent_find.cpp)
 add_rocprim_benchmark(benchmark_device_batch_memcpy.cpp)
 add_rocprim_benchmark(benchmark_device_binary_search.cpp)
 add_rocprim_benchmark(benchmark_device_find_first_of.cpp)
+add_rocprim_benchmark(benchmark_device_find_end.cpp)
 add_rocprim_benchmark(benchmark_device_histogram.cpp)
 add_rocprim_benchmark(benchmark_device_merge.cpp)
 add_rocprim_benchmark(benchmark_device_merge_sort.cpp)
@@ -156,7 +144,9 @@ add_rocprim_benchmark(benchmark_device_run_length_encode.cpp)
 add_rocprim_benchmark(benchmark_device_scan.cpp)
 add_rocprim_benchmark(benchmark_device_scan_deterministic.cpp)
 add_rocprim_benchmark(benchmark_device_scan_by_key.cpp)
+add_rocprim_benchmark(benchmark_device_search.cpp)
 add_rocprim_benchmark(benchmark_device_scan_by_key_deterministic.cpp)
+add_rocprim_benchmark(benchmark_device_search_n.cpp)
 add_rocprim_benchmark(benchmark_device_select.cpp)
 add_rocprim_benchmark(benchmark_device_segmented_radix_sort_keys.cpp)
 add_rocprim_benchmark(benchmark_device_segmented_radix_sort_pairs.cpp)

benchmark/ConfigAutotuneSettings.cmake

@@ -33,6 +33,10 @@ function(read_config_autotune_settings file list_across_names list_across output
     set(list_across "${TUNING_TYPES};\
 true;false true;32 64 128 256 512 1024" PARENT_SCOPE)
     set(output_pattern_suffix "@DataType@_@Left@_@InPlace@_@BlockSize@" PARENT_SCOPE)
+  elseif(file STREQUAL "benchmark_device_adjacent_find")
+    set(list_across_names "InputType;BlockSize" PARENT_SCOPE)
+    set(list_across "${TUNING_TYPES};64 128 256 512 1024" PARENT_SCOPE)
+    set(output_pattern_suffix "@InputType@_@BlockSize@" PARENT_SCOPE)
   elseif(file STREQUAL "benchmark_device_histogram")
     set(list_across_names "DataType;BlockSize" PARENT_SCOPE)
     set(list_across "${TUNING_TYPES};64 128 256" PARENT_SCOPE)
@@ -115,5 +119,9 @@ DataType;BlockSize;" PARENT_SCOPE)
     set(list_across_names "DataType;BlockSize" PARENT_SCOPE)
     set(list_across "${LIMITED_TUNING_TYPES};32 64 128 256 512 1024" PARENT_SCOPE)
     set(output_pattern_suffix "@DataType@_@BlockSize@" PARENT_SCOPE)
+  elseif(file STREQUAL "benchmark_device_merge")
+    set(list_across_names "KeyType;ValueType;BlockSize" PARENT_SCOPE)
+    set(list_across "${TUNING_TYPES};rocprim::empty_type ${LIMITED_TUNING_TYPES};32 64 128 256 512 1024" PARENT_SCOPE)
+    set(output_pattern_suffix "@KeyType@_@ValueType@_@BlockSize@" PARENT_SCOPE)
   endif()
 endfunction()