Develop stream 2025-03-05 (#697)

* Resolve "Extract common utilities for benchmarks and tests"

* Add missing error check

* implement dynamic dispatch for batch_copy and batch_memcpy

* Reduce merge inplace memory footprint

* Resolve "block disconinuity test values are converted to long long before comparing"

* Resolve "Remove debug builds and combine benchmark and tests in Windows CI"

* Resolve "Broadcast algorithm in warp_scan MI300X / MI325X"

* Resolve "hipGetDeviceProperties is slow"

* Resolve "Move device_ptr to common utilities"

* Resolve "Use existing enums in thread_load_kernel() and thread_store_kernel()"

* Resolve "Generate testing data using rocRAND"

* Resolve "Explicitly declared values in enum for thread_load and thread_store modifiers"

* Reduce lds conflicts shared histogram

* Reset the user's active device on error

* Resolve "Remove comment in test_device_histogram"

* Resolve "Calculating the individual positions of each item dynamically, rather than storing them in registers"

* Resolve "Implementation Proposal Virtual Shared Memory"

* Resolve "fix precision check per operations in scan and reduce"

* Use number_of_runs in loops for tests

* Resolve "update CHANGELOG.md to include changes by new tests data generation"

* Added warning for deprecation of c++14

* Remove duplicate include

---------

Co-authored-by: Yung-Sheng Tu <[email protected]>
Co-authored-by: Sander Bos <[email protected]>
Co-authored-by: Bence Parajdi <[email protected]>
Co-authored-by: Nara Prasetya <[email protected]>
Co-authored-by: Robin Voetter <[email protected]>
Co-authored-by: Enrico Degregori <[email protected]>
Co-authored-by: Milo Lurati <[email protected]>
Co-authored-by: Borys Petrov <[email protected]>
Co-authored-by: Mátyás Aradi <[email protected]>

Author: 10 people

.gitlab-ci.yml

@@ -188,6 +188,7 @@ build:cmake-minimum-apt:
       -D CMAKE_CXX_FLAGS="-Wall -Wextra -Werror $EXTRA_CMAKE_CXX_FLAGS"
       -D CMAKE_BUILD_TYPE="$BUILD_TYPE"
       -D BUILD_$BUILD_TARGET=ON
+      -D WITH_ROCRAND=ON
       -D CMAKE_C_COMPILER_LAUNCHER=phc_sccache_c
       -D CMAKE_CXX_COMPILER_LAUNCHER=phc_sccache_cxx
       -D BUILD_EXAMPLE=ON
@@ -206,6 +207,7 @@ build:cmake-minimum-apt:
       - $BUILD_DIR/gtest/
       - $BUILD_DIR/test/CTestTestfile.cmake
       - $BUILD_DIR/test/rocprim/CTestTestfile.cmake
+      - $BUILD_DIR/deps/rocrand/
       - $BUILD_DIR/test/rocprim/test_*
       - $BUILD_DIR/test/rocprim/libtest_*
       - $BUILD_DIR/test/test_*
@@ -271,14 +273,14 @@ build:windows:
     - .deps:visual-studio-devshell
   parallel:
     matrix:
-      - BUILD_TYPE: [Debug, Release]
-        BUILD_TARGET: [BENCHMARK, TEST]
+      - BUILD_TYPE: Release
   script:
     - mkdir -p $CI_PROJECT_DIR/build
     - cmake -G Ninja
       -S $CI_PROJECT_DIR
       -B $CI_PROJECT_DIR/build
-      -D BUILD_$BUILD_TARGET=ON
+      -D BUILD_TEST=ON
+      -D BUILD_BENCHMARK=ON
       -D AMDGPU_TARGETS=$GPU_TARGET
       -D CMAKE_CXX_COMPILER:PATH="${env:HIP_PATH}\bin\clang++.exe"
       -D CMAKE_PREFIX_PATH:PATH="${env:HIP_PATH}"
@@ -427,7 +429,6 @@ test-windows-release:
       parallel:
         matrix:
           - BUILD_TYPE: Release
-            BUILD_TARGET: TEST
 
 .test-package:
   script:

CHANGELOG.md

@@ -17,6 +17,7 @@ Full documentation for rocPRIM is available at [https://rocm.docs.amd.com/projec
 * 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.
+* Added support for building tests with device-side random data generation, making them finish faster. This requires rocRAND, and is enabled with the `WITH_ROCRAND=ON` build flag.
 
 ### Changed
 

CMakeLists.txt

@@ -68,6 +68,7 @@ include(CMakeDependentOption)
 # Disables building tests, benchmarks, examples
 option(ONLY_INSTALL "Only install" OFF)
 cmake_dependent_option(BUILD_TEST "Build tests (requires googletest)" OFF "NOT ONLY_INSTALL" OFF)
+option(WITH_ROCRAND "Build tests with device-side data generation(requires rocRAND)" OFF)
 cmake_dependent_option(BUILD_BENCHMARK "Build benchmarks" OFF "NOT ONLY_INSTALL" OFF)
 cmake_dependent_option(BUILD_EXAMPLE "Build examples" OFF "NOT ONLY_INSTALL" OFF)
 option(BUILD_NAIVE_BENCHMARK "Build naive benchmarks" OFF)
@@ -131,6 +132,7 @@ set(BUILD_SHARED_LIBS OFF) # don't build client dependencies as shared
 
 # Get dependencies (required here to get rocm-cmake)
 include(cmake/Dependencies.cmake)
+
 # Use target ID syntax if supported for GPU_TARGETS
 if(USE_HIPCXX)
   enable_language(HIP)

CONTRIBUTING.md

@@ -41,6 +41,8 @@ and performs well across a variety of input types and sizes. More specifically:
 - Tests and benchmarks must be instantiated with all supported data types.
   - If the algorithm uses multiple data types (for instance, if it uses different types for input and output), a selected and representative few combinations should be tested instead of the full combination matrix.
 
+Any utility needed by the tests **and** benchmarks must be added to the appropriate header within the `common` folder. Non-common utilities may be hosted in the corresponding headers from the `test` or `benchmark` folders. For a more detailed description of the cases to be considered for adding new utilities, please check [common](/common/README.md).
+
 We also employ automated testing and benchmarking via checks that are run when a pull request is created.
 These checks:
 - test all algorithms for correctness across a variety of input configurations (eg. types, sizes, etc.)

benchmark/ConfigAutotuneSettings.cmake

@@ -29,10 +29,10 @@ set(LIMITED_TUNING_TYPES "rocprim::int128_t int64_t int short int8_t")
 
 function(read_config_autotune_settings file list_across_names list_across output_pattern_suffix)
   if(file STREQUAL "benchmark_device_adjacent_difference")
-    set(list_across_names "DataType;Left;InPlace;BlockSize" PARENT_SCOPE)
+    set(list_across_names "DataType;Left;Aliasing;BlockSize" PARENT_SCOPE)
     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)
+true;no_alias in_place;32 64 128 256 512 1024" PARENT_SCOPE)
+    set(output_pattern_suffix "@DataType@_@Left@_@Aliasing@_@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)

benchmark/benchmark_block_exchange.cpp

@@ -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
@@ -24,6 +24,8 @@
 #include "benchmark_utils.hpp"
 #include "cmdparser.hpp"
 
+#include "../common/utils_custom_type.hpp"
+
 // Google Benchmark
 #include <benchmark/benchmark.h>
 
@@ -312,8 +314,8 @@ void add_benchmarks(const std::string&                            name,
                     const managed_seed&                           seed,
                     hipStream_t                                   stream)
 {
-    using custom_float2  = custom_type<float, float>;
-    using custom_double2 = custom_type<double, double>;
+    using custom_float2  = common::custom_type<float, float>;
+    using custom_double2 = common::custom_type<double, double>;
 
     std::vector<benchmark::internal::Benchmark*> bs = {BENCHMARK_TYPE(int, 256),
                                                        BENCHMARK_TYPE(int8_t, 256),

benchmark/benchmark_block_radix_rank.cpp

@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2022-2024 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2022-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
@@ -24,6 +24,8 @@
 // CmdParser
 #include "cmdparser.hpp"
 
+#include "../common/utils_data_generation.hpp"
+
 // Google Benchmark
 #include <benchmark/benchmark.h>
 
@@ -110,8 +112,8 @@ void run_benchmark(benchmark::State&   state,
     const unsigned int     size            = items_per_block * grid_size;
 
     std::vector<T> input = get_random_data<T>(size,
-                                              generate_limits<T>::min(),
-                                              generate_limits<T>::max(),
+                                              common::generate_limits<T>::min(),
+                                              common::generate_limits<T>::max(),
                                               seed.get_0());
 
     T*            d_input;

benchmark/benchmark_block_radix_sort.cpp

@@ -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
@@ -24,6 +24,9 @@
 // CmdParser
 #include "cmdparser.hpp"
 
+#include "../common/utils_custom_type.hpp"
+#include "../common/utils_data_generation.hpp"
+
 // Google Benchmark
 #include <benchmark/benchmark.h>
 
@@ -54,7 +57,7 @@ enum class benchmark_kinds
 };
 
 template<typename T>
-using select_decomposer_t = std::conditional_t<is_custom_type<T>::value,
+using select_decomposer_t = std::conditional_t<common::is_custom_type<T>::value,
                                                custom_type_decomposer<T>,
                                                rocprim::identity_decomposer>;
 
@@ -140,8 +143,8 @@ void run_benchmark(benchmark::State&   state,
     const auto     size = items_per_block * ((N + items_per_block - 1) / items_per_block);
 
     std::vector<T> input = get_random_data<T>(size,
-                                              generate_limits<T>::min(),
-                                              generate_limits<T>::max(),
+                                              common::generate_limits<T>::min(),
+                                              common::generate_limits<T>::max(),
                                               seed.get_0());
 
     T* d_input;
@@ -230,7 +233,7 @@ void add_benchmarks(benchmark_kinds                               benchmark_kind
                     const managed_seed&                           seed,
                     hipStream_t                                   stream)
 {
-    using custom_int_type = custom_type<int, int>;
+    using custom_int_type = common::custom_type<int, int>;
 
     std::vector<benchmark::internal::Benchmark*> bs = {BENCHMARK_TYPE(int, 64, 3),
                                                        BENCHMARK_TYPE(int, 512, 3),

benchmark/benchmark_block_reduce.cpp

@@ -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
@@ -24,6 +24,8 @@
 // CmdParser
 #include "cmdparser.hpp"
 
+#include "../common/utils_custom_type.hpp"
+
 // Google Benchmark
 #include <benchmark/benchmark.h>
 
@@ -170,8 +172,8 @@ void add_benchmarks(std::vector<benchmark::internal::Benchmark*>& benchmarks,
                     hipStream_t                                   stream,
                     size_t                                        bytes)
 {
-    using custom_float2  = custom_type<float, float>;
-    using custom_double2 = custom_type<double, double>;
+    using custom_float2  = common::custom_type<float, float>;
+    using custom_double2 = common::custom_type<double, double>;
 
     std::vector<benchmark::internal::Benchmark*> new_benchmarks
         = {// When block size is less than or equal to warp size

benchmark/benchmark_block_run_length_decode.cpp

@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2021-2024 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2021-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
@@ -23,6 +23,8 @@
 #include "benchmark_utils.hpp"
 #include "cmdparser.hpp"
 
+#include "../common/utils_data_generation.hpp"
+
 #include <benchmark/benchmark.h>
 
 #include <rocprim/block/block_load_func.hpp>
@@ -115,10 +117,10 @@ void run_benchmark(benchmark::State&   state,
 
     engine_type prng(seed.get_0());
     using ItemDistribution = std::conditional_t<rocprim::is_integral<ItemT>::value,
-                                                uniform_int_distribution<ItemT>,
+                                                common::uniform_int_distribution<ItemT>,
                                                 std::uniform_real_distribution<ItemT>>;
-    ItemDistribution                  run_item_dist(0, 100);
-    uniform_int_distribution<OffsetT> run_length_dist(MinRunLength, MaxRunLength);
+    ItemDistribution                          run_item_dist(0, 100);
+    common::uniform_int_distribution<OffsetT> run_length_dist(MinRunLength, MaxRunLength);
 
     for(size_t i = 0; i < num_runs; ++i)
     {

benchmark/benchmark_block_scan.cpp

@@ -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
@@ -24,6 +24,8 @@
 // CmdParser
 #include "cmdparser.hpp"
 
+#include "../common/utils_custom_type.hpp"
+
 // Google Benchmark
 #include <benchmark/benchmark.h>
 
@@ -205,8 +207,8 @@ void add_benchmarks(std::vector<benchmark::internal::Benchmark*>& benchmarks,
                     hipStream_t                                   stream,
                     size_t                                        bytes)
 {
-    using custom_float2  = custom_type<float, float>;
-    using custom_double2 = custom_type<double, double>;
+    using custom_float2  = common::custom_type<float, float>;
+    using custom_double2 = common::custom_type<double, double>;
 
     std::vector<benchmark::internal::Benchmark*> new_benchmarks
         = {// When block size is less than or equal to warp size

benchmark/benchmark_block_sort.parallel.hpp

@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2019-2024 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2019-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
@@ -25,6 +25,8 @@
 
 #include "benchmark_utils.hpp"
 
+#include "../common/utils_data_generation.hpp"
+
 // Google Benchmark
 #include <benchmark/benchmark.h>
 
@@ -243,10 +245,11 @@ struct block_sort_benchmark : public config_autotune_interface
 
         const auto size = items_per_block * ((N + items_per_block - 1) / items_per_block);
 
-        std::vector<KeyType> input = get_random_data<KeyType>(size,
-                                                              generate_limits<KeyType>::min(),
-                                                              generate_limits<KeyType>::max(),
-                                                              seed.get_0());
+        std::vector<KeyType> input
+            = get_random_data<KeyType>(size,
+                                       common::generate_limits<KeyType>::min(),
+                                       common::generate_limits<KeyType>::max(),
+                                       seed.get_0());
 
         KeyType* d_input;
         KeyType* d_output;

benchmark/benchmark_device_adjacent_difference.cpp

@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2022-2024 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2022-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
@@ -23,6 +23,11 @@
 #include "benchmark_device_adjacent_difference.parallel.hpp"
 #include "benchmark_utils.hpp"
 
+#ifndef BENCHMARK_CONFIG_TUNING
+    #include "../common/device_adjacent_difference.hpp"
+    #include "../common/utils_custom_type.hpp"
+#endif
+
 // Google Benchmark
 #include <benchmark/benchmark.h>
 
@@ -56,10 +61,10 @@ constexpr std::size_t DEFAULT_BYTES = 1024LL * 1024LL * 1024LL * 2LL;
 
 // clang-format off
 #define CREATE_BENCHMARKS(T)          \
-    CREATE_BENCHMARK(T, true,  false) \
-    CREATE_BENCHMARK(T, true,  true)  \
-    CREATE_BENCHMARK(T, false, false) \
-    CREATE_BENCHMARK(T, false, true)
+    CREATE_BENCHMARK(T, true,  common::api_variant::no_alias) \
+    CREATE_BENCHMARK(T, true,  common::api_variant::in_place)  \
+    CREATE_BENCHMARK(T, false, common::api_variant::no_alias) \
+    CREATE_BENCHMARK(T, false, common::api_variant::in_place)
 // clang-format on
 
 int main(int argc, char* argv[])
@@ -112,8 +117,8 @@ int main(int argc, char* argv[])
                                                         seed,
                                                         stream);
 #else // BENCHMARK_CONFIG_TUNING
-    using custom_float2  = custom_type<float, float>;
-    using custom_double2 = custom_type<double, double>;
+    using custom_float2  = common::custom_type<float, float>;
+    using custom_double2 = common::custom_type<double, double>;
     // Add benchmarks
     CREATE_BENCHMARKS(int)
     CREATE_BENCHMARKS(std::int64_t)

benchmark/benchmark_device_adjacent_difference.parallel.cpp.in

@@ -1,6 +1,6 @@
 // MIT License
 //
-// Copyright (c) 2022-2024 Advanced Micro Devices, Inc. All rights reserved.
+// Copyright (c) 2022-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
@@ -20,18 +20,20 @@
 // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
 // SOFTWARE.
 
-#include "benchmark_utils.hpp"
 #include "benchmark_device_adjacent_difference.parallel.hpp"
+#include "benchmark_utils.hpp"
+
+#include "../common/device_adjacent_difference.hpp"
 
 #include <rocprim/types.hpp>
 
 #include <stdint.h>
 
 namespace {
     auto benchmarks = config_autotune_register::create_bulk(
-        device_adjacent_difference_benchmark_generator< 
-        @DataType@, 
+        device_adjacent_difference_benchmark_generator<
+        @DataType@,
         @BlockSize@,
-        @Left@, 
-        @InPlace@>::create);
+        @Left@,
+        common::api_variant::@Aliasing@>::create);
 }