Develop stream 2024-06-26

.clang-format

@@ -41,7 +41,6 @@ AllowShortLoopsOnASingleLine: false
 AlwaysBreakAfterReturnType: None
 AlwaysBreakBeforeMultilineStrings: false
 AlwaysBreakTemplateDeclarations: Yes
-AttributeMacros: ['ROCPRIM_DEVICE', 'ROCPRIM_HOST', 'ROCPRIM_HOST_DEVICE', 'ROCPRIM_SHARED_MEMORY', 'ROCPRIM_KERNEL', 'ROCPRIM_INLINE']
 BinPackArguments: false
 BinPackParameters: false
 BitFieldColonSpacing: Both
@@ -135,4 +134,40 @@ SpacesInConditionalStatement: false
 SpacesInContainerLiterals: true
 SpacesInParentheses: false
 SpacesInSquareBrackets: false
+
+AttributeMacros:
+    - __host__
+    - __device__
+    - __global__
+    - __forceinline__
+    - __shared__
+    - __launch_bounds__
+    - ROCPRIM_DEVICE
+    - ROCPRIM_HOST
+    - ROCPRIM_HOST_DEVICE
+    - ROCPRIM_SHARED_MEMORY
+    - ROCPRIM_KERNEL
+    - ROCPRIM_INLINE
+    - ROCPRIM_FORCE_INLINE
+    - ROCPRIM_LAUNCH_BOUNDS
+
+# Trick clang into thinking that our C-style attributes are C++-style attributes
+# Make sure that the sizes line up for linebreaks etc
+Macros:
+    - __host__=[[host]]
+    - __device__=[[device]]
+    - __global__=[[global]]
+    - __forceinline__=[[forceinline]]
+    - __shared__=[[shared]]
+    - __launch_bounds__(x)=[[launch_bounds(x)]]
+    - __attribute__(x)=[[attribute(x)]]
+    - ROCPRIM_DEVICE=[[DEVICE____]]
+    - ROCPRIM_HOST=[[HOST____]]
+    - ROCPRIM_HOST_DEVICE=[[HOST_DEVICE____]]
+    - ROCPRIM_SHARED_MEMORY=[[SHARED_MEMORY____]]
+    - ROCPRIM_KERNEL=[[KERNEL____]]
+    - ROCPRIM_INLINE=[[INLINE____]]
+    - ROCPRIM_FORCE_INLINE=[FORCE_INLINE____]]
+    - ROCPRIM_LAUNCH_BOUNDS(x)=[[launch_bounds(x)____]]
+BreakAfterAttributes: Always
 ---

.gitlab-ci.yml

@@ -27,6 +27,7 @@ include:
       - /defaults.yaml
       - /deps-cmake.yaml
       - /deps-docs.yaml
+      - /deps-format.yaml
       - /deps-rocm.yaml
       - /deps-vcpkg.yaml
       - /deps-windows.yaml
@@ -46,20 +47,7 @@ variables:
 
 clang-format:
   extends:
-    - .deps:rocm
-  stage: lint
-  needs: []
-  tags:
-    - build
-  variables:
-    CLANG_FORMAT: "/opt/rocm/llvm/bin/clang-format"
-    GIT_CLANG_FORMAT: "/opt/rocm/llvm/bin/git-clang-format"
-  rules:
-    - if: '$CI_PIPELINE_SOURCE == "merge_request_event"'
-  script:
-    - cd $CI_PROJECT_DIR
-    - git config --global --add safe.directory $CI_PROJECT_DIR
-    - scripts/code-format/check-format.sh $CI_MERGE_REQUEST_DIFF_BASE_SHA --binary "$CLANG_FORMAT"
+    - .lint:clang-format
 
 copyright-date:
   extends:
@@ -162,15 +150,22 @@ build:cmake-minimum-apt:
   extends:
     - .gpus:rocm-gpus
     - .rules:build
+  variables:
+    EXTRA_CMAKE_CXX_FLAGS: ""
   script:
     - mkdir -p $BUILD_DIR
     - cd $BUILD_DIR
+    - | # 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"
+        fi
     - cmake
       -G Ninja
       -D CMAKE_CXX_COMPILER="$AMDCLANG"
-      -D CMAKE_CXX_FLAGS="-Wall -Wextra -Werror"
+      -D CMAKE_CXX_FLAGS="-Wall -Wextra -Werror $EXTRA_CMAKE_CXX_FLAGS"
       -D CMAKE_BUILD_TYPE="$BUILD_TYPE"
-      -D BUILD_$BUILD_TARGET=ON 
+      -D BUILD_$BUILD_TARGET=ON
       -D BUILD_EXAMPLE=ON
       -D GPU_TARGETS=$GPU_TARGETS
       -D AMDGPU_TEST_TARGETS=$GPU_TARGETS
@@ -197,9 +192,11 @@ build:cmake-latest:
   extends:
     - .cmake-latest
     - .build:common
-  variables:
-    BUILD_TYPE: Release
-    BUILD_TARGET: TEST
+  parallel:
+    # Debug builds disabled due to excessive build times for debug test builds
+    matrix:
+      - BUILD_TYPE: Release
+        BUILD_TARGET: [BENCHMARK, TEST]
 
 build:cmake-minimum:
   needs: []
@@ -246,20 +243,19 @@ build:windows:
     - .deps:visual-studio-devshell
   parallel:
     matrix:
-      # Debug is disabled due to extensive link times, tracked in issue 679.
-      - BUILD_TYPE: [Release]
+      - BUILD_TYPE: [Debug, Release]
         BUILD_TARGET: [BENCHMARK, TEST]
   script:
     - mkdir -p $CI_PROJECT_DIR/build
-    - cmake -G Ninja 
-      -S $CI_PROJECT_DIR 
-      -B $CI_PROJECT_DIR/build 
-      -D BUILD_$BUILD_TARGET=ON 
+    - cmake -G Ninja
+      -S $CI_PROJECT_DIR
+      -B $CI_PROJECT_DIR/build
+      -D BUILD_$BUILD_TARGET=ON
       -D GPU_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" 
-    - cmake --build "$CI_PROJECT_DIR/build" 
+      -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"
+    - cmake --build "$CI_PROJECT_DIR/build"
   artifacts:
     paths:
       - $CI_PROJECT_DIR/build/test/test_*
@@ -281,6 +277,10 @@ autotune:build:
     - .cmake-minimum
     - .gpus:rocm-gpus
     - .rules:benchmark
+  before_script:
+    - !reference [".cmake-minimum", before_script]
+    - $SUDO_CMD apt-get update -qq
+    - $SUDO_CMD apt-get install -qq -y zstd
   variables:
     BENCHMARK_TARGETS: benchmark_config_tuning
   script:
@@ -301,22 +301,13 @@ autotune:build:
       -D GPU_TARGETS=$GPU_TARGETS
     - cmake --build . --target $BENCHMARK_TARGETS
     - 'rm -rf $BUILD_DIR/benchmark/benchmark*.parallel'
-    # remove benchmark executables if their size together is too large for gitlab ci to handle
-    - |
-      total_size_bytes=0
-      while read -r file_size; do
-        total_size_bytes=$((total_size_bytes + file_size))
-      done < <(stat --format="%s" benchmark/benchmark*)
-      total_size_gib="$(numfmt --round=down --to-unit=Gi "$total_size_bytes")"
-      if [ "$total_size_gib" -ge 3 ]; then
-        printf "Total size: %s (%d bytes) > 3GiB, skipping benchmark executables from the artifact.\n" \
-          "$(numfmt --to=iec-i "$total_size_bytes")" "$total_size_bytes"
-        rm benchmark/benchmark*
-      fi
+    # The autotune benchmarks get very large, above GitLabs upload limit. Fortunately they compress well.
+    # We'll put them all in a single archive to compress them to a few hundred MB.
+    - find benchmark -type f -executable -print0 | tar -I zstd -cvf benchmarks.tar.zstd --null -T -
 
   artifacts:
     paths:
-      - $BUILD_DIR/benchmark/benchmark*
+      - $BUILD_DIR/benchmarks.tar.zstd
       - $BUILD_DIR/.ninja_log
       - $BUILD_DIR/deps/googlebenchmark/
     expire_in: 1 week
@@ -360,7 +351,7 @@ test:
     - cd $CI_PROJECT_DIR/build
     - ctest --output-on-failure
 
-# Disabled due to extensive link times. 
+# Disabled due to extensive link times.
 # This is tracked in issue 679
 # test-windows-debug:
 #   extends:
@@ -535,9 +526,13 @@ autotune:execute-tuning:
   artifacts:
     paths:
       - ${AUTOTUNE_RESULT_DIR}/*.json
+  before_script:
+    - !reference [".cmake-minimum", before_script]
+    - $SUDO_CMD apt-get update -qq
+    - $SUDO_CMD apt-get install -qq -y zstd
   script:
-    - >
-      cd "${CI_PROJECT_DIR}"
+    - cd "${CI_PROJECT_DIR}"
+    - tar -I zstd -xvf "${BUILD_DIR}/benchmarks.tar.zstd" -C "${BUILD_DIR}/"
     - |
       if [ ! -d "${BUILD_DIR}/benchmark" ]; then
           echo "There are no benchmark executables. Run the build job with a BUILD_TARGET."
@@ -553,6 +548,7 @@ autotune:execute-tuning:
       --benchmark_filter_regex="${AUTOTUNE_ALGORITHM_REGEX}"
       --size="${AUTOTUNE_SIZE}"
       --trials="${AUTOTUNE_TRIALS}"
+      --seed=82589933
 
 autotune:generate-config:
   image: python:3.10.5-buster

.gitlab/run_benchmarks.py

@@ -1,6 +1,6 @@
 #!/usr/bin/env python3
 
-# Copyright (c) 2022-2023 Advanced Micro Devices, Inc. All rights reserved.
+# Copyright (c) 2022-2024 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
@@ -28,7 +28,7 @@
 import subprocess
 import sys
 
-BenchmarkContext = namedtuple('BenchmarkContext', ['gpu_architecture', 'benchmark_output_dir', 'benchmark_dir', 'benchmark_filename_regex', 'benchmark_filter_regex', 'size', 'trials'])
+BenchmarkContext = namedtuple('BenchmarkContext', ['gpu_architecture', 'benchmark_output_dir', 'benchmark_dir', 'benchmark_filename_regex', 'benchmark_filter_regex', 'size', 'trials', 'seed'])
 
 def run_benchmarks(benchmark_context):
     def is_benchmark_executable(filename):
@@ -61,6 +61,8 @@ def is_benchmark_executable(filename):
             args += ['--size', benchmark_context.size]
         if benchmark_context.trials:
             args += ['--trials', benchmark_context.trials]
+        if benchmark_context.seed:
+            args += ['--seed', benchmark_context.seed]
         try:
             subprocess.check_call(args)
         except subprocess.CalledProcessError as error:
@@ -97,6 +99,10 @@ def main():
         help='Controls the number of trial iterations for each benchmark case',
         default='',
         required=False)
+    parser.add_argument('--seed',
+        help='Controls the seed for random number generation for each benchmark case',
+        default='',
+        required=False)
 
     args = parser.parse_args()
 
@@ -107,7 +113,8 @@ def main():
         args.benchmark_filename_regex,
         args.benchmark_filter_regex,
         args.size,
-        args.trials)
+        args.trials,
+        args.seed)
 
     benchmark_run_successful = run_benchmarks(benchmark_context)
 

CHANGELOG.md

@@ -5,11 +5,36 @@ Documentation for rocPRIM is available at
 
 ## Unreleased rocPRIM-3.3.0 for ROCm 6.3.0
 
-### Fixes
+### Additions
+
+* Option `--seed` to benchmarks to specify a seed for the generation of random inputs. The default behavior is to keep using a random seed per benchmark measurement.
+* Added configuration autotuning to device partition (`rocprim::partition`, `rocprim::partition_two_way`, and `rocprim::partition_three_way`), device select (`rocprim::select`, `rocprim::unique`, and `rocprim::unique_by_key`), and device reduce by key (`rocprim::reduce_by_key`) for improved performance on selected architectures.
+* Added `rocprim::uninitialized_array` which provides uninitialized storage in local memory for user-defined types.
+* Added large segment support for `rocprim:segmented_reduce`.
+* Added a parallel `nth_element` device function similar to `std::nth_element`, this function rearranges elements smaller than the n-th before and bigger than the n-th after the n-th element.
+* Added deterministic (bitwise reproducible) algorithm variants `rocprim::deterministic_inclusive_scan`, `rocprim::deterministic_exclusive_scan`, `rocprim::deterministic_inclusive_scan_by_key`, `rocprim::deterministic_exclusive_scan_by_key`, and `rocprim::deterministic_reduce_by_key`. These provide run-to-run stable results with non-associative operators such as float operations, at the cost of reduced performance.
+* Added a parallel `partial_sort` and `partial_sort_copy` device function similar to `std::partial_sort` and `std::partial_sort_copy`, these functions rearranges elements such that the elements are the same as a sorted list up to and including the middle index.
+
+### Changes
 
+* Modified the input size in device adjacent difference benchmarks. Observed performance with these benchmarks might be different.
+* Changed the default seed for `device_benchmark_segmented_reduce`.
+
+### Fixes
 
 * Fixed an issue where while running rtest.py on windows and passing in an absolute path to `--install_dir` causes a `FileNotFound` error.
+* rocPRIM functions are no longer forcefully inlined on Windows, significantly reducing the build
+  time in debug builds.
+* `block_load`, `block_store`, `block_shuffle`, `block_exchange` and `warp_exchange` now use placement `new` instead of copy
+  assignment (`operator=`) when writing to local memory. This fixes the behavior of custom types with non-trivial copy assignments.
+* Fixed a bug in the generation of input data for benchmarks, which caused incorrect performance to be reported in specific cases. It may affect the reported performance for one-byte types (`uint8_t` and `int8_t`) and instantiations of `custom_type`. Specifically, device binary search, device histogram, device merge and warp sort are affected.
+* Fixed a bug for `rocprim::merge_path_search` where using `unsigned` offsets would output wrong results.
+* Fixed a bug for `rocprim::thread_load` and `rocprim::thread_store` where `float` and `double` were not casted to the correct type resulting in wrong results.
+* Fix tests failing when compiling with `-D_GLIBCXX_ASSERTIONS=ON`.
+
+### Deprecations
 
+* `rocprim::thread_load` and `rocprim::thread_store`, use dereference instead. Not all of those functions are available on every device architecture, and their usage can hurt performance, because inline assembly inhibits optimizations.
 
 ## Unreleased rocPRIM-3.2.0 for ROCm 6.2.0
 
@@ -28,10 +53,10 @@ Documentation for rocPRIM is available at
 * New `rocprim::batch_copy` function added. Similar to `rocprim::batch_memcpy`, but copies by element, not with memcpy.
 * Added more test cases, to better cover supported data types.
 * Updated some tests to work with supported data types.
-* An optional `decomposer` argument for all member functions of `rocprim::block_radix_sort` and all functions of `device_radix_sort`. 
+* An optional `decomposer` argument for all member functions of `rocprim::block_radix_sort` and all functions of `device_radix_sort`.
   To sort keys of an user-defined type, a decomposer functor should be passed. The decomposer should produce a `rocprim::tuple`
   of references to arithmetic types from the key.
-* New `rocprim::predicate_iterator` which acts as a proxy for an underlying iterator based on a predicate. 
+* New `rocprim::predicate_iterator` which acts as a proxy for an underlying iterator based on a predicate.
   It iterates over proxies that holds the references to the underlying values, but only allow reading and writing if the predicate is `true`.
   It can be instantiated with:
   * `rocprim::make_predicate_iterator`
@@ -43,6 +68,7 @@ Documentation for rocPRIM is available at
 
 * Improved the performance of `warp_sort_shuffle` and `block_sort_bitonic`.
 * Created an optimized version of the `warp_exchange` functions `blocked_to_striped_shuffle` and `striped_to_blocked_shuffle` when the warpsize is equal to the items per thread.
+* Improved the performance of `device_transform`.
 
 ### Fixes
 

CMakeLists.txt

@@ -45,6 +45,7 @@ endif()
 # Build options
 option(BUILD_TEST "Build tests (requires googletest)" OFF)
 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)
@@ -130,7 +131,7 @@ if(USE_HIP_CPU)
 endif()
 
 # Setup VERSION
-set(VERSION_STRING "3.2.0")
+set(VERSION_STRING "3.3.0")
 rocm_setup_version(VERSION ${VERSION_STRING})
 
 # Print configuration summary