From 4df9baec03cfb06ad22fd1bb20b82f1346321836 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Fri, 17 Jul 2026 11:38:10 -0400 Subject: [PATCH 01/12] feat(rocthrust): grafting pr changes --- projects/rocthrust/CHANGELOG.md | 5 ++ projects/rocthrust/CMakeLists.txt | 31 ++++++---- projects/rocthrust/CMakePresets.json | 56 +++++++++++++++++ projects/rocthrust/README.md | 8 ++- projects/rocthrust/cmake/Benchmarks.cmake | 61 ------------------- projects/rocthrust/cmake/Dependencies.cmake | 61 ++----------------- projects/rocthrust/cmake/Summary.cmake | 5 +- projects/rocthrust/cmake/VerifyCompiler.cmake | 8 +-- projects/rocthrust/test/CMakeLists.txt | 10 +-- projects/rocthrust/testing/CMakeLists.txt | 4 +- projects/rocthrust/testing/allocator.cu | 1 - .../testing/dependencies_aware_policies.cu | 9 ++- projects/rocthrust/testing/unittest/meta.h | 5 -- 13 files changed, 115 insertions(+), 149 deletions(-) create mode 100644 projects/rocthrust/CMakePresets.json delete mode 100644 projects/rocthrust/cmake/Benchmarks.cmake diff --git a/projects/rocthrust/CHANGELOG.md b/projects/rocthrust/CHANGELOG.md index f77dc955cca4..9b424da832d4 100644 --- a/projects/rocthrust/CHANGELOG.md +++ b/projects/rocthrust/CHANGELOG.md @@ -14,6 +14,11 @@ Documentation for rocThrust available at * CCCL 2.8.x compatibility is deprecated. hipCUB and rocThrust will be brought forward to CCCL 3.0.x compatibility in an upcoming version. +### Changed + +* Benchmarking now uses primbench for its benchmarks instead of Google Benchmark. + * See `shared/primbench/README.md` for its documentation. + ## Since last release ROCm 7.12 ### Resolved issues diff --git a/projects/rocthrust/CMakeLists.txt b/projects/rocthrust/CMakeLists.txt index 467bb97f7549..989c177cd6b3 100644 --- a/projects/rocthrust/CMakeLists.txt +++ b/projects/rocthrust/CMakeLists.txt @@ -68,9 +68,11 @@ option(ROCTHRUST_USE_LIBHIPCXX "Build rocThrust using libhipcxx backend, if avai option(ROCTHRUST_USE_LIBCUDACXX "Build rocThrust using libcudacxx backend, if available" ON) cmake_dependent_option(ENABLE_UPSTREAM_TESTS "Enable upstream (thrust) tests" ON BUILD_TEST OFF) cmake_dependent_option(USE_SYSTEM_LIB "Use existing system ROCm library installation when building tests" OFF BUILD_TEST OFF) -option(EXTERNAL_DEPS_FORCE_DOWNLOAD "Force download of non-ROCm dependencies (eg. Google Test, Google Benchmark)" OFF) +option(EXTERNAL_DEPS_FORCE_DOWNLOAD "Force download of non-ROCm dependencies (e.g., Google Test)" OFF) option(LINK_HIP_DEVICE_LIBS "Links against HIP device libraries. Required for device-side acceleration." ON) +option(GRAFT_THRUST_ONTO_BINARIES "!!! TODO !!!" OFF) + if((BUILD_TEST OR BUILD_BENCHMARK OR BUILD_EXAMPLE) AND NOT (${LINK_HIP_DEVICE_LIBS})) message(FATAL_ERROR "rocThrust tests, benchmarks, and examples must be built with LINK_HIP_DEVICE_LIBS=ON since they require device acceleration.") endif() @@ -178,7 +180,7 @@ check_option(FETCH_METHOD_OPTIONS ROCRAND_FETCH_METHOD) include(cmake/Dependencies.cmake) # Verify that supported compilers are used -if (NOT WIN32) +if (NOT WIN32 OR NOT GRAFT_THRUST_ONTO_BINARIES) include(cmake/VerifyCompiler.cmake) endif() @@ -191,11 +193,13 @@ set(THRUST_HOST_SYSTEM CPP CACHE STRING "The host backend to target.") check_option(THRUST_HOST_SYSTEM_OPTIONS THRUST_HOST_SYSTEM) add_compile_definitions(THRUST_HOST_SYSTEM=THRUST_HOST_SYSTEM_${THRUST_HOST_SYSTEM}) -# Note CUDA backend is currently not supported, so disallow the option for now. -set(THRUST_DEVICE_SYSTEM_OPTIONS OMP TBB CPP HIP) -set(THRUST_DEVICE_SYSTEM HIP CACHE STRING "The host backend to target.") -check_option(THRUST_DEVICE_SYSTEM_OPTIONS THRUST_DEVICE_SYSTEM) -add_compile_definitions(THRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_${THRUST_DEVICE_SYSTEM}) +# If we're not grafting existing CUDA onto binaries, we must dissallow CUDA backend. +if (NOT GRAFT_THRUST_ONTO_BINARIES) + set(THRUST_DEVICE_SYSTEM_OPTIONS OMP TBB CPP HIP) + set(THRUST_DEVICE_SYSTEM HIP CACHE STRING "The host backend to target.") + check_option(THRUST_DEVICE_SYSTEM_OPTIONS THRUST_DEVICE_SYSTEM) + add_compile_definitions(THRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_${THRUST_DEVICE_SYSTEM}) +endif() set(COMPILE_OPTIONS -Wall -Wextra) if(NOT DISABLE_WERROR) @@ -230,8 +234,13 @@ endif() rocm_setup_version(VERSION ${VERSION_STRING}) math(EXPR rocthrust_VERSION_NUMBER "${rocthrust_VERSION_MAJOR} * 100000 + ${rocthrust_VERSION_MINOR} * 100 + ${rocthrust_VERSION_PATCH}") -# Thrust (with HIP backend) -add_subdirectory(thrust) +# Thrust (header-only) library +if(GRAFT_THRUST_ONTO_BINARIES) + # If we're grafting CCCL we don't need the actual rocThrust library! + find_package(CCCL CONFIG REQUIRED) +else() + add_subdirectory(thrust) +endif() if(BUILD_TEST OR BUILD_BENCHMARK OR BUILD_HIPSTDPAR_TEST) rocm_package_setup_component(clients) @@ -242,7 +251,7 @@ if(CODE_COVERAGE) endif() # Tests -if(BUILD_TEST AND USE_SYSTEM_LIB) +if(BUILD_TEST AND USE_SYSTEM_LIB AND NOT GRAFT_THRUST_ONTO_BINARIES) find_package(rocprim REQUIRED CONFIG PATHS "/opt/rocm/rocprim") if (${rocprim_VERSION} VERSION_LESS ${MIN_ROCPRIM_PACKAGE_VERSION}) message(WARNING "The installed rocprim version, ${rocprim_VERSION}, is less than the minimum required version ${MIN_ROCPRIM_PACKAGE_VERSION}. Building tests with USE_SYSTEM_LIB=ON may not work properly.") @@ -261,7 +270,7 @@ if(BUILD_TEST OR BUILD_HIPSTDPAR_TEST) # Used by tests to define their RESOURCE_GROUPS set(DEFAULT_GPU_RESOURCE "gpus") - + # We still want the testing to be compiled to catch some errors #TODO: Get testing folder working with HIP on Windows if (NOT WIN32 AND BUILD_TEST) diff --git a/projects/rocthrust/CMakePresets.json b/projects/rocthrust/CMakePresets.json new file mode 100644 index 000000000000..a0d9507d3ce5 --- /dev/null +++ b/projects/rocthrust/CMakePresets.json @@ -0,0 +1,56 @@ +{ + "version": 3, + "cmakeMinimumRequired": { + "major": 3, + "minor": 28, + "patch": 0 + }, + "configurePresets": [ + { + "name": "hip", + "hidden": true, + "generator": "Ninja", + "binaryDir": "${sourceDir}/build/${presetName}", + "cacheVariables": { + "CMAKE_CXX_STANDARD": "17", + "CMAKE_BUILD_TYPE": "Release", + "CMAKE_EXPORT_COMPILE_COMMANDS": true + } + }, + { + "inherits": "hip", + "name": "hip-amd", + "cacheVariables": { + "HIP_PLATFORM": "amd" + } + }, + { + "inherits": "hip-amd", + "name": "hip-amd-all", + "cacheVariables": { + "BUILD_BENCHMARK": true, + "BUILD_TEST": true + } + }, + { + "inherits": "hip-amd-all", + "name": "hip-amd-dev", + "cacheVariables": { + "ROCRAND_FETCH_METHOD": "MONOREPO" + } + }, + { + "inherits": "hip", + "name": "hip-nv-dev", + "cacheVariables": { + "BUILD_BENCHMARK": true, + "CMAKE_HIP_FLAGS": "--extended-lambda -w", + "GRAFT_THRUST_ONTO_BINARIES": true, + "HIP_PLATFORM": "nvidia", + "ROCRAND_FETCH_METHOD": "MONOREPO", + "THRUST_DEVICE_SYSTEM": "CUDA", + "USE_HIPCXX": true + } + } + ] +} diff --git a/projects/rocthrust/README.md b/projects/rocthrust/README.md index d1c78ff4b854..3c53fa00d200 100644 --- a/projects/rocthrust/README.md +++ b/projects/rocthrust/README.md @@ -30,6 +30,8 @@ Optional: * This is automatically downloaded and built by the CMake script * [doxygen](https://www.doxygen.nl/) * Required for building the documentation +* [AMD SMI](https://github.com/ROCm/amdsmi) + * Required only for benchmarks. Building benchmarks is off by default. For ROCm hardware requirements, refer to: @@ -73,7 +75,7 @@ mkdir build; cd build # BUILD_BENCHMARK - OFF by default, # ROCPRIM_FETCH_METHOD - PACKAGE is the default, see below for a description of available options # ROCRAND_FETCH_METHOD - PACKAGE is the default, see below for a description of available options -# 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) # BUILD_ADDRESS_SANITIZER - OFF by default, builds with clang address sanitizer enabled. # RNG_SEED_COUNT - 0 by default, controls non-repeatable random dataset count # PRNG_SEEDS - 1 by default, reproducible seeds to generate random data @@ -106,7 +108,7 @@ rocThrust can (optionally) automatically fetch a number of dependencies for you Alternatively, it can seach for an existing installation on your system. The following cmake build options control how dependencies are located. -- `EXTERNAL_DEPS_FORCE_DOWNLOAD` (default: `OFF`) - when set to `ON`, non-ROCm dependencies (Google Test, Google Benchmark) will always be downloaded, even if they are already installed ony your system. When set to `OFF`, rocThrust first searches for existing installations of the dependencies on your system, and only downloads them if they cannot be found. +- `EXTERNAL_DEPS_FORCE_DOWNLOAD` (default: `OFF`) - when set to `ON`, non-ROCm dependencies (Google Test) will always be downloaded, even if they are already installed ony your system. When set to `OFF`, rocThrust first searches for existing installations of the dependencies on your system, and only downloads them if they cannot be found. - `ROCPRIM_FETCH_METHOD` (default: `PACKAGE`) - controls the way that the rocPRIM dependency is fetched. This option must be set to one of: - `PACKAGE` - searches for an existing installation of the dependency. If it is not found, rocThrust will fall back to using the `DOWNLOAD` setting (below). @@ -292,6 +294,8 @@ make -j4 ## Running benchmarks +rocThrust 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. + ```sh # Go to rocThrust build directory cd projects/rocthrust; cd build diff --git a/projects/rocthrust/cmake/Benchmarks.cmake b/projects/rocthrust/cmake/Benchmarks.cmake deleted file mode 100644 index 8437add5e4ef..000000000000 --- a/projects/rocthrust/cmake/Benchmarks.cmake +++ /dev/null @@ -1,61 +0,0 @@ -# ######################################################################## -# Copyright 2024-2025 Advanced Micro Devices, Inc. -# ######################################################################## - -# ########################### -# rocThrust benchmarks -# ########################### - -# Common functionality for configuring rocThrust's benchmarks - -# Registers a .cu as C++ rocThrust benchmark -function(add_thrust_benchmark BENCHMARK_NAME BENCHMARK_SOURCE NOT_INTERNAL) - set(BENCHMARK_TARGET "benchmark_thrust_${BENCHMARK_NAME}") - if(USE_HIPCXX) - set_source_files_properties(${BENCHMARK_SOURCE} - PROPERTIES - LANGUAGE HIP - ) - else() - set_source_files_properties(${BENCHMARK_SOURCE} - PROPERTIES - LANGUAGE CXX - ) - endif() - add_executable(${BENCHMARK_TARGET} ${BENCHMARK_SOURCE}) - - target_compile_options(${BENCHMARK_TARGET} PRIVATE ${COMPILE_OPTIONS}) - target_link_libraries(${BENCHMARK_TARGET} - PRIVATE - rocthrust - roc::rocprim_hip - ) - # Internal benchmark does not use Google Benchmark nor rocRAND. - # This can be omited when that benchmark is removed. - if(NOT_INTERNAL) - target_link_libraries(${BENCHMARK_TARGET} - PRIVATE - roc::rocrand - benchmark::benchmark - ) - endif() - foreach(gpu_target ${GPU_TARGETS}) - target_link_libraries(${BENCHMARK_TARGET} - INTERFACE - --cuda-gpu-arch=${gpu_target} - ) - endforeach() - - # Separate normal from internal benchmarks - if(NOT_INTERNAL) - set(OUTPUT_DIR "${CMAKE_BINARY_DIR}/benchmark/") - else() - set(OUTPUT_DIR "${CMAKE_BINARY_DIR}/benchmark/internal/") - endif() - - set_target_properties(${BENCHMARK_TARGET} - PROPERTIES - RUNTIME_OUTPUT_DIRECTORY "${OUTPUT_DIR}" - ) - rocm_install(TARGETS ${BENCHMARK_TARGET} COMPONENT benchmarks) -endfunction() diff --git a/projects/rocthrust/cmake/Dependencies.cmake b/projects/rocthrust/cmake/Dependencies.cmake index 6044443c641c..edebc17fd859 100644 --- a/projects/rocthrust/cmake/Dependencies.cmake +++ b/projects/rocthrust/cmake/Dependencies.cmake @@ -1,5 +1,5 @@ # ######################################################################## -# Copyright 2019-2025 Advanced Micro Devices, Inc. +# Copyright 2019-2026 Advanced Micro Devices, Inc. # ######################################################################## # ########################### @@ -244,7 +244,7 @@ function(fetch_dep method repo_name repo_path download_branch) endif() endfunction() -if(${LINK_HIP_DEVICE_LIBS}) +if(${LINK_HIP_DEVICE_LIBS} AND NOT GRAFT_THRUST_ONTO_BINARIES) fetch_dep(ROCPRIM_FETCH_METHOD rocprim ROCPRIM_PATH ROCM_DEP_RELEASE_BRANCH) if(${ROCPRIM_FETCH_METHOD} STREQUAL "DOWNLOAD" OR ${ROCPRIM_FETCH_METHOD} STREQUAL "MONOREPO") @@ -254,7 +254,7 @@ if(${LINK_HIP_DEVICE_LIBS}) prim SOURCE_DIR ${ROCPRIM_PATH} INSTALL_DIR ${CMAKE_CURRENT_BINARY_DIR}/deps/rocprim - CMAKE_ARGS -DBUILD_TEST=OFF -DCMAKE_INSTALL_PREFIX= -DCMAKE_PREFIX_PATH=/opt/rocm + CMAKE_ARGS -DBUILD_TEST=OFF -DBUILD_BENCHMARK=OFF -DBUILD_EXAMPLE=OFF -DCMAKE_INSTALL_PREFIX= -DCMAKE_PREFIX_PATH=/opt/rocm LOG_CONFIGURE TRUE LOG_BUILD TRUE LOG_INSTALL TRUE @@ -377,57 +377,6 @@ endif() # Benchmark dependencies if(BUILD_BENCHMARK) - set(BENCHMARK_VERSION 1.9.5) - if(NOT EXTERNAL_DEPS_FORCE_DOWNLOAD) - # Google Benchmark (https://github.com/google/benchmark.git) - find_package(benchmark ${BENCHMARK_VERSION} QUIET) - else() - message(STATUS "Force installing Google Benchmark.") - endif() - - if(NOT benchmark_FOUND) - message(STATUS "Google Benchmark not found or force download Google Benchmark on. Downloading and building Google Benchmark.") - if(CMAKE_CONFIGURATION_TYPES) - message(FATAL_ERROR "DownloadProject.cmake doesn't support multi-configuration generators.") - endif() - set(GOOGLEBENCHMARK_ROOT ${CMAKE_CURRENT_BINARY_DIR}/deps/googlebenchmark CACHE PATH "") - if(NOT (CMAKE_CXX_COMPILER_ID STREQUAL "GNU")) - if(WIN32) - get_filename_component(CXX_DIRNAME ${CMAKE_CXX_COMPILER} DIRECTORY) - set(COMPILER_OVERRIDE "-DCMAKE_CXX_COMPILER=${CXX_DIRNAME}/clang++.exe") - else() - set(COMPILER_OVERRIDE "-DCMAKE_CXX_COMPILER=g++") - endif() - endif() - - message(STATUS "Google Benchmark not found. Fetching...") - option(BENCHMARK_ENABLE_TESTING "Enable testing of the benchmark library." OFF) - option(BENCHMARK_ENABLE_INSTALL "Enable installation of benchmark." OFF) - FetchContent_Declare( - googlebench - GIT_REPOSITORY https://github.com/google/benchmark.git - GIT_TAG v${BENCHMARK_VERSION} - ) - set(HAVE_STD_REGEX ON) - set(RUN_HAVE_STD_REGEX 1) - set(_ROCTHRUST_DISABLE_ROCM_CHECKS TRUE) - FetchContent_MakeAvailable(googlebench) - set(_ROCTHRUST_DISABLE_ROCM_CHECKS FALSE) - # Clang on Windows throws the following warnings with Googlebenchmark v1.9.5 (along with Werror): - # googlebench-src/src/string_util.cc:158:34: error: format string is not a string literal [-Werror,-Wformat-nonliteral] - if(CMAKE_CXX_COMPILER_ID STREQUAL "Clang" AND WIN32) - if(TARGET benchmark) - target_compile_options(benchmark PRIVATE -Wno-format-nonliteral -Wno-missing-format-attribute -Wno-unused-command-line-argument) - endif() - if(TARGET benchmark_main) - target_compile_options(benchmark_main PRIVATE -Wno-format-nonliteral -Wno-missing-format-attribute -Wno-unused-command-line-argument) - endif() - if(NOT TARGET benchmark::benchmark) - add_library(benchmark::benchmark ALIAS benchmark) - endif() - endif() - endif() - # rocRAND (https://github.com/ROCm/rocm-libraries) fetch_dep(ROCRAND_FETCH_METHOD rocrand ROCRAND_PATH ROCM_DEP_RELEASE_BRANCH) @@ -443,12 +392,12 @@ if(BUILD_BENCHMARK) if(CMAKE_CXX_COMPILER_LAUNCHER) set(EXTRA_CMAKE_ARGS "${EXTRA_CMAKE_ARGS} -DCMAKE_CXX_COMPILER_LAUNCHER=${CMAKE_CXX_COMPILER_LAUNCHER}") endif() - + FetchContent_Declare( rocrand SOURCE_DIR ${ROCRAND_PATH} INSTALL_DIR ${CMAKE_CURRENT_BINARY_DIR}/deps/rocrand - CMAKE_ARGS -DCMAKE_INSTALL_PREFIX= -DCMAKE_PREFIX_PATH=/opt/rocm ${EXTRA_CMAKE_ARGS} + CMAKE_ARGS -DBUILD_BENCHMARK=OFF -DBUILD_TEST=OFF -DBUILD_EXAMPLE=OFF -DCMAKE_INSTALL_PREFIX= -DCMAKE_PREFIX_PATH=/opt/rocm ${EXTRA_CMAKE_ARGS} LOG_CONFIGURE TRUE LOG_BUILD TRUE LOG_INSTALL TRUE diff --git a/projects/rocthrust/cmake/Summary.cmake b/projects/rocthrust/cmake/Summary.cmake index f38d1b6f461e..cae85aa9cc4b 100644 --- a/projects/rocthrust/cmake/Summary.cmake +++ b/projects/rocthrust/cmake/Summary.cmake @@ -83,9 +83,10 @@ endif() message(STATUS " Install prefix : ${CMAKE_INSTALL_PREFIX}") if(HIP_COMPILER STREQUAL "clang") if(USE_HIPCXX) - message(STATUS " Device targets : ${CMAKE_HIP_ARCHITECTURES}") + message(STATUS " CMAKE_HIP_ARCHITECTURES: ${CMAKE_HIP_ARCHITECTURES}") + message(STATUS " HIP_ARCHITECTURES: ${HIP_ARCHITECTURES}") else() - message(STATUS " Device targets : ${GPU_TARGETS}") + message(STATUS " GPU_TARGETS: ${GPU_TARGETS}") endif() endif() message(STATUS "") diff --git a/projects/rocthrust/cmake/VerifyCompiler.cmake b/projects/rocthrust/cmake/VerifyCompiler.cmake index 3c0fc81200b8..4b183e574682 100644 --- a/projects/rocthrust/cmake/VerifyCompiler.cmake +++ b/projects/rocthrust/cmake/VerifyCompiler.cmake @@ -25,9 +25,9 @@ if(${LINK_HIP_DEVICE_LIBS}) find_package(hip REQUIRED CONFIG PATHS /opt/rocm) endif() -if(HIP_COMPILER STREQUAL "nvcc") - message(FATAL_ERROR "rocThrust does not support the CUDA backend.") -endif() +# if(HIP_COMPILER STREQUAL "nvcc") +# message(FATAL_ERROR "rocThrust does not support the CUDA backend.") +# endif() if(${LINK_HIP_DEVICE_LIBS}) # When building for HIP, make sure we have a hip-aware clang. @@ -41,7 +41,5 @@ if(${LINK_HIP_DEVICE_LIBS}) message(FATAL_ERROR "When LINK_HIP_DEVICE_LIBS is set to 'ON', then 'hipcc' or a HIP-aware Clang must be used as the C++ compiler.") endif() endif() - else() - message(FATAL_ERROR "When LINK_HIP_DEVICE_LIBS is set to 'ON', HIP_COMPILER must be `clang` (AMD ROCm platform)") endif() endif() diff --git a/projects/rocthrust/test/CMakeLists.txt b/projects/rocthrust/test/CMakeLists.txt index 031be6512739..d2c526bb6b67 100644 --- a/projects/rocthrust/test/CMakeLists.txt +++ b/projects/rocthrust/test/CMakeLists.txt @@ -64,11 +64,13 @@ function(add_rocthrust_test TEST) ${sqlite_local_SOURCE_DIR} ) - if(USE_SYSTEM_LIB) - target_link_libraries(${TEST_TARGET} PRIVATE roc::rocthrust) + if(GRAFT_THRUST_ONTO_BINARIES) + target_link_libraries(testing_common INTERFACE CCCL::CCCL) + elseif(USE_SYSTEM_LIB) + target_link_libraries(testing_common INTERFACE roc::rocthrust) else() - target_link_libraries(${TEST_TARGET} - PRIVATE + target_link_libraries(testing_common + INTERFACE rocthrust roc::rocprim_hip ) diff --git a/projects/rocthrust/testing/CMakeLists.txt b/projects/rocthrust/testing/CMakeLists.txt index 0ddb6bf03e91..9c35606c4bf9 100644 --- a/projects/rocthrust/testing/CMakeLists.txt +++ b/projects/rocthrust/testing/CMakeLists.txt @@ -10,7 +10,9 @@ target_include_directories(testing_common $ ) -if(USE_SYSTEM_LIB) +if(GRAFT_THRUST_ONTO_BINARIES) + target_link_libraries(testing_common INTERFACE CCCL::CCCL) +elseif(USE_SYSTEM_LIB) target_link_libraries(testing_common INTERFACE roc::rocthrust) else() target_link_libraries(testing_common diff --git a/projects/rocthrust/testing/allocator.cu b/projects/rocthrust/testing/allocator.cu index 43332eccdd25..d8caa554b57c 100644 --- a/projects/rocthrust/testing/allocator.cu +++ b/projects/rocthrust/testing/allocator.cu @@ -17,7 +17,6 @@ #include -#include #include #include diff --git a/projects/rocthrust/testing/dependencies_aware_policies.cu b/projects/rocthrust/testing/dependencies_aware_policies.cu index 3153ef825aa2..ff9278159b1c 100644 --- a/projects/rocthrust/testing/dependencies_aware_policies.cu +++ b/projects/rocthrust/testing/dependencies_aware_policies.cu @@ -22,10 +22,14 @@ THRUST_SUPPRESS_DEPRECATED_PUSH #include #include -#include #include #include +// Skip testing the HIP policy when not targeting HIP. +#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_HIP +#include +#endif + #include #if !_THRUST_HAS_DEVICE_SYSTEM_STD @@ -118,6 +122,8 @@ using cpp_par_info = policy_info; using tbb_par_info = policy_info; +// Skip testing the HIP policy when not targeting HIP. +#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_HIP using hip_par_info = policy_info; SimpleUnitTest> TestDependencyAttachmentInstance; +#endif THRUST_SUPPRESS_DEPRECATED_POP diff --git a/projects/rocthrust/testing/unittest/meta.h b/projects/rocthrust/testing/unittest/meta.h index 71cc4d47a856..3c42c02689bc 100644 --- a/projects/rocthrust/testing/unittest/meta.h +++ b/projects/rocthrust/testing/unittest/meta.h @@ -23,11 +23,6 @@ #pragma once -// This header `` is added only for suppressing the -// warning generated by deprecated API `thrust::device_malloc_allocator`. Once the -// API is removed, this header should also be removed. -#include - namespace unittest { From 32abfbd267945be1d7b78af082eb64dc85257e18 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Fri, 17 Jul 2026 11:39:41 -0400 Subject: [PATCH 02/12] feat(rocthrust): begin primbench integeration --- projects/rocthrust/.gitlab/run_benchmarks.py | 355 ++++++-- projects/rocthrust/benchmark/CMakeLists.txt | 179 ++-- projects/rocthrust/benchmark/bench_utils.hpp | 254 ++++++ .../benchmark/bench_utils/bench_utils.hpp | 734 ---------------- .../benchmark/bench_utils/cmdparser.hpp | 657 -------------- .../benchmark/bench_utils/common/types.hpp | 97 -- .../benchmark/bench_utils/custom_reporter.hpp | 826 ------------------ .../{bench_utils => }/generation_utils.hpp | 232 +---- 8 files changed, 700 insertions(+), 2634 deletions(-) create mode 100644 projects/rocthrust/benchmark/bench_utils.hpp delete mode 100644 projects/rocthrust/benchmark/bench_utils/bench_utils.hpp delete mode 100644 projects/rocthrust/benchmark/bench_utils/cmdparser.hpp delete mode 100644 projects/rocthrust/benchmark/bench_utils/common/types.hpp delete mode 100644 projects/rocthrust/benchmark/bench_utils/custom_reporter.hpp rename projects/rocthrust/benchmark/{bench_utils => }/generation_utils.hpp (71%) diff --git a/projects/rocthrust/.gitlab/run_benchmarks.py b/projects/rocthrust/.gitlab/run_benchmarks.py index 6a94c25c615a..317ae7515d50 100755 --- a/projects/rocthrust/.gitlab/run_benchmarks.py +++ b/projects/rocthrust/.gitlab/run_benchmarks.py @@ -1,6 +1,6 @@ #!/usr/bin/env python3 -# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. +# Copyright (c) 2024-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 @@ -21,100 +21,325 @@ # THE SOFTWARE. import argparse -from collections import namedtuple +import json import os import re import stat import subprocess import sys -import time -BenchmarkContext = namedtuple('BenchmarkContext', ['gpu_architecture', 'benchmark_output_dir', 'benchmark_dir', 'benchmark_filename_regex', 'benchmark_filter_regex', 'seed']) -def run_benchmarks(benchmark_context): +def run_benchmarks( + benchmark_executables_dir, + gpu_architecture, + json_out_dir, + csv_out_dir, + benchmark_executable_filter, + skip_gathered, + size, + hot, + seed, + filter, + dry, + min_gpu_ms_per_batch, + min_secs, + noise_timeout_secs, + batch_window_size, + noise_tolerance_percent, + min_gpu_temp, + max_gpu_temp, + max_warming_secs, + max_cooling_secs, + output_batches, + spaces_per_indent, + stream_blocking_timeout_secs, +): def is_benchmark_executable(filename): - if not re.match(benchmark_context.benchmark_filename_regex, filename): + if not re.search(benchmark_executable_filter, filename): return False - path = os.path.join(benchmark_context.benchmark_dir, filename) + + path = os.path.join(benchmark_executables_dir, filename) st_mode = os.stat(path).st_mode - # we are not interested in permissions, just whether there is any execution flag set - # and it is a regular file (S_IFREG) - return (st_mode & (stat.S_IXUSR | stat.S_IXGRP | stat.S_IXOTH)) and (st_mode & stat.S_IFREG) + # Return True when any execution flag is set, AND it is a regular file. + # Doesn't check permissions. + return (st_mode & (stat.S_IXUSR | stat.S_IXGRP | stat.S_IXOTH)) and ( + st_mode & stat.S_IFREG + ) + + def should_skip(json_out_path): + if not skip_gathered: + return False + + try: + with open(json_out_path) as f: + json.load(f) + except (FileNotFoundError, json.JSONDecodeError): + return False + + return True + + # Create output directories if they don't exist + os.makedirs(json_out_dir, exist_ok=True) + if csv_out_dir: + os.makedirs(csv_out_dir, exist_ok=True) + + benchmark_names = [ + name + for name in os.listdir(benchmark_executables_dir) + if is_benchmark_executable(name) + ] + print( + "The following benchmarks will be run:\n{}".format("\n".join(benchmark_names)), + file=sys.stderr, + flush=True, + ) success = True - benchmark_names = [name for name in os.listdir(benchmark_context.benchmark_dir) if is_benchmark_executable(name)] - benchmark_names.sort() - print('The following benchmarks will be ran:\n{}'.format('\n'.join(benchmark_names)), file=sys.stderr, flush=True) + for benchmark_name in benchmark_names: - results_json_name = f'{benchmark_name}_{benchmark_context.gpu_architecture}.json' - - benchmark_path = os.path.join(benchmark_context.benchmark_dir, benchmark_name) - results_json_path = os.path.join(benchmark_context.benchmark_output_dir, results_json_name) - args = [ - benchmark_path, - '--name_format', - 'json', - '--benchmark_out_format=json', - f'--benchmark_out={results_json_path}', - f'--benchmark_filter={benchmark_context.benchmark_filter_regex}' - ] - if benchmark_context.seed: - args += ['--seed', benchmark_context.seed] + results_json_name = f"{benchmark_name}_{gpu_architecture}.json" + json_out_path = os.path.join(json_out_dir, results_json_name) + + if should_skip(json_out_path): + print( + f"Skipping {benchmark_name}, because its results have already been gathered at {json_out_path}", + file=sys.stderr, + flush=True, + ) + continue + + benchmark_path = os.path.join(benchmark_executables_dir, benchmark_name) + + args = [benchmark_path] + args += ["--json-out", json_out_path] + + if size: + args += ["--size", size] + if hot: + args += ["--hot"] + if seed: + args += ["--seed", seed] + if csv_out_dir: + results_csv_name = f"{benchmark_name}_{gpu_architecture}.csv" + csv_out_path = os.path.join(csv_out_dir, results_csv_name) + args += ["--csv-out", csv_out_path] + if filter: + args += ["--filter", filter] + if dry: + args += ["--dry"] + if min_gpu_ms_per_batch: + args += ["--min-gpu-ms-per-batch", min_gpu_ms_per_batch] + if min_secs: + args += ["--min-secs", min_secs] + if noise_timeout_secs: + args += ["--noise-timeout-secs", noise_timeout_secs] + if batch_window_size: + args += ["--batch-window-size", batch_window_size] + if noise_tolerance_percent: + args += ["--noise-tolerance-percent", noise_tolerance_percent] + if min_gpu_temp: + args += ["--min-gpu-temp", min_gpu_temp] + if max_gpu_temp: + args += ["--max-gpu-temp", max_gpu_temp] + if max_warming_secs: + args += ["--max-warming-secs", max_warming_secs] + if max_cooling_secs: + args += ["--max-cooling-secs", max_cooling_secs] + if output_batches: + args += ["--output-batches"] + if spaces_per_indent: + args += ["--spaces-per-indent", spaces_per_indent] + if stream_blocking_timeout_secs: + args += ["--stream-blocking-timeout-secs", stream_blocking_timeout_secs] + try: - start_time = time.time() subprocess.check_call(args) - end_time = time.time() - duration = end_time - start_time - - print(f'Benchmark {benchmark_name} took {duration:.3f} seconds to run', file=sys.stderr, flush=True) except subprocess.CalledProcessError as error: - print(f'Could not run benchmark at {benchmark_path}. Error: "{error}"', file=sys.stderr, flush=True) + print( + f'Could not run benchmark at {benchmark_path}. Error: "{error}"', + file=sys.stderr, + flush=True, + ) success = False - return success + return success def main(): parser = argparse.ArgumentParser() - parser.add_argument('--benchmark_dir', - help='The local directory that contains the benchmark executables', - required=True) - parser.add_argument('--benchmark_gpu_architecture', - help='The architecture of the currently enabled GPU', - required=True) - parser.add_argument('--benchmark_output_dir', - help='The directory to write the benchmarks to', - required=True) - parser.add_argument('--benchmark_filename_regex', - help='Regular expression that controls the list of benchmark executables to run', - default=r'^benchmark', - required=False) - parser.add_argument('--benchmark_filter_regex', - help='Regular expression that controls the list of benchmarks to run in each benchmark executable', - 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() + # Python arguments + parser.add_argument( + "--benchmark-executables-dir", + help="Benchmark executables directory", + required=True, + ) + parser.add_argument( + "--gpu-architecture", + help="Architecture of the currently enabled GPU", + required=True, + ) + parser.add_argument( + "--json-out-dir", + help="Directory to write JSON result files to, creating parent directories if necessary", + required=True, + ) + parser.add_argument( + "--csv-out-dir", + help="Directory to write CSV result files to, creating parent directories if necessary", + required=False, + ) + parser.add_argument( + "--benchmark-executable-filter", + help="Regular expression that controls the list of benchmark executables to run", + default=r"^benchmark", + required=False, + ) + parser.add_argument( + "--skip-gathered", + help="Skip running benchmarks whose JSON file has already been gathered", + default=False, + action="store_true", + required=False, + ) - benchmark_context = BenchmarkContext( - args.benchmark_gpu_architecture, - args.benchmark_output_dir, - args.benchmark_dir, - args.benchmark_filename_regex, - args.benchmark_filter_regex, - args.seed) + # primbench arguments + parser.add_argument( + "--size", + help="Input size. Benchmarks decide what this represents, but it is commonly the number of bytes or items", + default="", + required=False, + ) + parser.add_argument( + "--hot", + help="Skip clearing the GPU cache between batch iterations", + default=False, + action="store_true", + required=False, + ) + parser.add_argument( + "--seed", help="Seed used for input generation", default="", required=False + ) + parser.add_argument( + "--filter", + help="Regex filter of specialization names to benchmark", + default="", + required=False, + ) + parser.add_argument( + "--dry", + help="Perform a dry run. The benchmark setup is still run, and JSON and CSV files are still output, but state.run() immediately returns", + default=False, + action="store_true", + required=False, + ) + parser.add_argument( + "--min-gpu-ms-per-batch", + help="Minimum duration of a batch in milliseconds (GPU time)", + default="", + required=False, + ) + parser.add_argument( + "--min-secs", + help="Minimum total benchmark duration in seconds (wall time)", + default="", + required=False, + ) + parser.add_argument( + "--noise-timeout-secs", + help="Maximum total benchmark duration in seconds before timing out a noisy run (wall time)", + default="", + required=False, + ) + parser.add_argument( + "--batch-window-size", + help="Number of batch times used in the noise (coefficient of variation) window to decide early stopping", + default="", + required=False, + ) + parser.add_argument( + "--noise-tolerance-percent", + help="Noise tolerance of batch times in percent, used to determine early benchmark stopping", + default="", + required=False, + ) + parser.add_argument( + "--min-gpu-temp", + help="Minimum GPU temperature in °C. Too low slows benchmarks; too high increases noise", + default="", + required=False, + ) + parser.add_argument( + "--max-gpu-temp", + help="Maximum GPU temperature in °C. Too low slows benchmarks; too high increases noise", + default="", + required=False, + ) + parser.add_argument( + "--max-warming-secs", + help="Maximum seconds allowed for GPU warming before an error is thrown", + default="", + required=False, + ) + parser.add_argument( + "--max-cooling-secs", + help="Maximum seconds allowed for GPU cooling before an error is thrown", + default="", + required=False, + ) + parser.add_argument( + "--output-batches", + help="Output a `batches` array for each specialization, containing per-batch details", + default=False, + action="store_true", + required=False, + ) + parser.add_argument( + "--spaces-per-indent", + help="Number of spaces per indentation level in JSON output. Set to 0 for no indentation", + default="", + required=False, + ) + parser.add_argument( + "--stream-blocking-timeout-secs", + help="Maximum stream blocking duration in seconds before timing out. Stream is blocked while queueing kernel calls. Use `primbench::flags::sync` if kernel is synchronous", + default="", + required=False, + ) + + args = parser.parse_args() - benchmark_run_successful = run_benchmarks(benchmark_context) + benchmark_run_successful = run_benchmarks( + args.benchmark_executables_dir, + args.gpu_architecture, + args.json_out_dir, + args.csv_out_dir, + args.benchmark_executable_filter, + args.skip_gathered, + args.size, + args.hot, + args.seed, + args.filter, + args.dry, + args.min_gpu_ms_per_batch, + args.min_secs, + args.noise_timeout_secs, + args.batch_window_size, + args.noise_tolerance_percent, + args.min_gpu_temp, + args.max_gpu_temp, + args.max_warming_secs, + args.max_cooling_secs, + args.output_batches, + args.spaces_per_indent, + args.stream_blocking_timeout_secs, + ) return benchmark_run_successful -if __name__ == '__main__': +if __name__ == "__main__": success = main() if success: exit(0) diff --git a/projects/rocthrust/benchmark/CMakeLists.txt b/projects/rocthrust/benchmark/CMakeLists.txt index 7eb9afd2549a..90e4a6f35a94 100644 --- a/projects/rocthrust/benchmark/CMakeLists.txt +++ b/projects/rocthrust/benchmark/CMakeLists.txt @@ -1,76 +1,131 @@ # ######################################################################## -# Copyright 2024 Advanced Micro Devices, Inc. +# Copyright 2024-2026 Advanced Micro Devices, Inc. # ######################################################################## -include(Benchmarks) +message(STATUS "Configuring benchmarks") -# Benchmarks directory in project's root -set(BENCHMARKS_ROOT "${CMAKE_CURRENT_LIST_DIR}") +# Force CMake to find Git immediately so ${GIT_EXECUTABLE} is populated +find_package(Git REQUIRED) -# Subdirectory of BENCHMARKS_ROOT containing all the .cu files grouped in -# subdirectories accordingly to the functionality benchmarked -set(BENCHMARKS_DIR "bench") +# 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() -# **************************************************************************** -# Functions -# **************************************************************************** +# 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 +) -# Gets all subdirs of benchmarks recursively -function(get_recursive_subdirs subdirs dir_prefix dir_name) - set(dirs) - file(GLOB_RECURSE contents - CONFIGURE_DEPENDS - LIST_DIRECTORIES ON - "${dir_prefix}/${dir_name}/*" - ) +function(add_thrust_benchmark BENCHMARK_SOURCE) + file(RELATIVE_PATH rel_path ${CMAKE_CURRENT_LIST_DIR}/bench ${CMAKE_CURRENT_LIST_DIR}/${BENCHMARK_SOURCE}) + string(REPLACE ".cu" "" bench_name ${rel_path}) + string(REPLACE "/" "_" bench_name ${bench_name}) - foreach(dir IN LISTS contents) - if(IS_DIRECTORY "${dir}") - list(APPEND dirs "${dir}") + set(BENCHMARK_TARGET benchmark_thrust_${bench_name}) + + if(USE_HIPCXX) + set_source_files_properties(${BENCHMARK_SOURCE} + PROPERTIES + LANGUAGE HIP + ) + else() + set_source_files_properties(${BENCHMARK_SOURCE} + PROPERTIES + LANGUAGE CXX + ) endif() - endforeach() - set(${subdirs} "${dirs}" PARENT_SCOPE) -endfunction() + add_executable(${BENCHMARK_TARGET} ${BENCHMARK_SOURCE}) + target_compile_options(${BENCHMARK_TARGET} PRIVATE ${COMPILE_OPTIONS}) -function(add_bench_dir bench_dir) - # Get algorithm name - get_filename_component(algo_name "${bench_dir}" NAME_WLE) - - # For scan, we also append the parent diretory name, as the algo_name - # will be exclusive/inclusive - get_filename_component(PARENT_DIR "${bench_dir}" DIRECTORY) - get_filename_component(PARENT_DIR_NAME "${PARENT_DIR}" NAME_WLE) - if(PARENT_DIR_NAME STREQUAL "scan") - set(algo_name "${PARENT_DIR_NAME}_${algo_name}") - endif() - - # Take all .cu (tests) inside the dir - file(GLOB bench_srcs CONFIGURE_DEPENDS "${bench_dir}/*.cu") - # Set benchmark prefix as its relative path to benchmarks separated - # by . instead of / - file(RELATIVE_PATH bench_prefix "${BENCHMARKS_ROOT}" "${bench_dir}") - file(TO_CMAKE_PATH "${bench_prefix}" bench_prefix) - string(REPLACE "/" "." bench_prefix "${bench_prefix}") - - # Add each benchmark as thrust benchmark - foreach(bench_src IN LISTS bench_srcs) - set(real_bench_src "${bench_src}") - # Get file name without directory nor last extension - get_filename_component(bench_name "${bench_src}" NAME_WLE) - add_thrust_benchmark("${algo_name}_${bench_name}" ${bench_src} ON) - endforeach() -endfunction() + if(GRAFT_THRUST_ONTO_BINARIES) + target_link_libraries(${BENCHMARK_TARGET} INTERFACE CCCL::CCCL) + else() + target_link_libraries(${BENCHMARK_TARGET} + PRIVATE + rocthrust + roc::rocprim_hip + ) + endif() + + target_link_libraries(${BENCHMARK_TARGET} PRIVATE roc::rocrand) -# **************************************************************************** -# Benchmarks -# **************************************************************************** -message (STATUS "Configuring benchmarks") + target_include_directories(${BENCHMARK_TARGET} PRIVATE + "${PROJECT_SOURCE_DIR}/benchmark" + "${PROJECT_SOURCE_DIR}/../../shared/primbench") -# Get all the subdirectories inside the bench directory -get_recursive_subdirs(subdirs ${BENCHMARKS_ROOT} ${BENCHMARKS_DIR}) + target_compile_definitions(${BENCHMARK_TARGET} PUBLIC BRANCH_NAME="${BRANCH_NAME}") + target_compile_definitions(${BENCHMARK_TARGET} PUBLIC COMMIT_HASH="${COMMIT_HASH}") + + if(WIN32) + target_compile_definitions(${BENCHMARK_TARGET} PRIVATE PRIMBENCH_NO_MONITORING) + elseif(HIP_COMPILER STREQUAL "nvcc") + target_link_libraries(${BENCHMARK_TARGET} PRIVATE nvidia-ml) + else() + target_link_libraries(${BENCHMARK_TARGET} PRIVATE amd_smi) + endif() + + foreach(gpu_target ${GPU_TARGETS}) + target_link_libraries(${BENCHMARK_TARGET} + INTERFACE + --cuda-gpu-arch=${gpu_target} + ) + endforeach() + + set_target_properties(${BENCHMARK_TARGET} + PROPERTIES + RUNTIME_OUTPUT_DIRECTORY ${CMAKE_BINARY_DIR}/benchmark/ + ) + rocm_install(TARGETS ${BENCHMARK_TARGET} COMPONENT benchmarks) +endfunction() -# Add benchmarks from each subdirectory present in bench -foreach(subdir IN LISTS subdirs) - add_bench_dir("${subdir}") -endforeach() +add_thrust_benchmark(bench/adjacent_difference/basic.cu) +add_thrust_benchmark(bench/adjacent_difference/custom.cu) +add_thrust_benchmark(bench/adjacent_difference/in_place.cu) +add_thrust_benchmark(bench/copy/basic.cu) +add_thrust_benchmark(bench/copy/if.cu) +add_thrust_benchmark(bench/equal/basic.cu) +add_thrust_benchmark(bench/fill/basic.cu) +add_thrust_benchmark(bench/for_each/basic.cu) +add_thrust_benchmark(bench/inner_product/basic.cu) +add_thrust_benchmark(bench/merge/basic.cu) +add_thrust_benchmark(bench/partition/basic.cu) +add_thrust_benchmark(bench/reduce/basic.cu) +add_thrust_benchmark(bench/reduce/by_key.cu) +add_thrust_benchmark(bench/scan/exclusive/by_key.cu) +add_thrust_benchmark(bench/scan/exclusive/max.cu) +add_thrust_benchmark(bench/scan/exclusive/sum.cu) +add_thrust_benchmark(bench/scan/inclusive/by_key.cu) +add_thrust_benchmark(bench/scan/inclusive/max.cu) +add_thrust_benchmark(bench/scan/inclusive/sum.cu) +add_thrust_benchmark(bench/set_operations/difference.cu) +add_thrust_benchmark(bench/set_operations/difference_by_key.cu) +add_thrust_benchmark(bench/set_operations/intersection.cu) +add_thrust_benchmark(bench/set_operations/intersection_by_key.cu) +add_thrust_benchmark(bench/set_operations/symmetric_difference.cu) +add_thrust_benchmark(bench/set_operations/symmetric_difference_by_key.cu) +add_thrust_benchmark(bench/set_operations/union.cu) +add_thrust_benchmark(bench/set_operations/union_by_key.cu) +add_thrust_benchmark(bench/shuffle/basic.cu) +add_thrust_benchmark(bench/sort/keys.cu) +add_thrust_benchmark(bench/sort/keys_custom.cu) +add_thrust_benchmark(bench/sort/pairs.cu) +add_thrust_benchmark(bench/sort/pairs_custom.cu) +add_thrust_benchmark(bench/tabulate/basic.cu) +add_thrust_benchmark(bench/transform/basic.cu) +add_thrust_benchmark(bench/transform_reduce/sum.cu) +add_thrust_benchmark(bench/unique/basic.cu) +add_thrust_benchmark(bench/unique/by_key.cu) +add_thrust_benchmark(bench/vectorized_search/basic.cu) +add_thrust_benchmark(bench/vectorized_search/lower_bound.cu) +add_thrust_benchmark(bench/vectorized_search/upper_bound.cu) diff --git a/projects/rocthrust/benchmark/bench_utils.hpp b/projects/rocthrust/benchmark/bench_utils.hpp new file mode 100644 index 000000000000..c49a21798753 --- /dev/null +++ b/projects/rocthrust/benchmark/bench_utils.hpp @@ -0,0 +1,254 @@ +// MIT License +// +// Copyright (c) 2024-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 +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in all +// copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +// SOFTWARE. + +#pragma once + +#include + +#include "primbench.hpp" + +#ifndef _MSC_VER +using int128_t = __int128_t; +using uint128_t = __uint128_t; +#endif + +#include "generation_utils.hpp" // IWYU pragma: export + +/// This allows running rocThrust benchmarks with CCCL Thrust. +#ifndef _THRUST_HAS_DEVICE_SYSTEM_STD +# define THRUST_HOST_DEVICE __host__ __device__ +# define THRUST_DEVICE __device__ + +# define _THRUST_LIBCXX_INCLUDE(LIB) +# define _THRUST_STD ::cuda::std +# define _THRUST_LIBCXX ::cuda +#endif // _THRUST_HAS_DEVICE_SYSTEM_STD + +#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_HIP +# include +#elif THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA +# include +#endif + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace bench_utils +{ + +#define HIP_CHECK(condition) \ + { \ + hipError_t error = condition; \ + if (error != hipSuccess) \ + { \ + std::cout << "HIP error: " << error << " line: " << __LINE__ << std::endl; \ + exit(error); \ + } \ + } + +struct less_t +{ + template + __host__ __device__ bool operator()(const T& lhs, const T& rhs) const + { + return lhs < rhs; + } +}; + +struct max_t +{ + template + __host__ __device__ T operator()(const T& lhs, const T& rhs) + { + less_t less{}; + return less(lhs, rhs) ? rhs : lhs; + } +}; + +/** + * This struct is used to reduce noise in benchmarks + */ +struct caching_allocator_t +{ + using value_type = char; + + caching_allocator_t() = default; + ~caching_allocator_t() + { + free_all(); + } + + char* allocate(std::ptrdiff_t num_bytes) + { + value_type* result{}; + auto free_block = free_blocks.find(num_bytes); + if (free_block != free_blocks.end()) + { + result = free_block->second; + free_blocks.erase(free_block); + } + else + { + HIP_CHECK(hipMalloc(&result, num_bytes)); + } + + allocated_blocks.emplace(result, num_bytes); + return result; + } + + void deallocate(value_type* ptr, size_t) + { + auto iter = allocated_blocks.find(ptr); + if (iter == allocated_blocks.end()) + { + throw std::runtime_error("Memory was not allocated by this allocator"); + } + + std::ptrdiff_t num_bytes = iter->second; + allocated_blocks.erase(iter); + free_blocks.emplace(num_bytes, ptr); + } + +private: + using FreeBlocksType = std::multimap; + using AllocatedBlocksType = std::map; + + FreeBlocksType free_blocks; + AllocatedBlocksType allocated_blocks; + + void free_all() + { + for (auto free_block : free_blocks) + { + HIP_CHECK(hipFree(free_block.second)); + } + + for (auto allocated_block : allocated_blocks) + { + HIP_CHECK(hipFree(allocated_block.first)); + } + } +}; + +class large_data +{ +public: + __host__ __device__ large_data() + { + data[0] = 0; + } + __host__ __device__ large_data(large_data const& val) + { + data[0] = val.data[0]; + } + __host__ __device__ large_data(int n) + { + data[0] = static_cast(n); + } + large_data& __host__ __device__ operator=(large_data const& val) + { + data[0] = val.data[0]; + return *this; + } + bool __host__ __device__ operator==(large_data const& val) const + { + return data[0] == val.data[0]; + } + large_data& __host__ __device__ operator++() + { + ++data[0]; + return *this; + } + __host__ __device__ operator int() const + { + return static_cast(data[0]); + } + + int8_t data[512]; +}; + +template +bool __host__ __device__ operator==(T const& lhs, large_data const& rhs) +{ + return static_cast(lhs).data[0] == rhs.data[0]; +} + +inline size_t total_global_mem() +{ + static const size_t mem = [] { + int device_id = 0; + HIP_CHECK(hipGetDevice(&device_id)); + hipDeviceProp_t prop{}; + HIP_CHECK(hipGetDeviceProperties(&prop, device_id)); + return prop.totalGlobalMem; + }(); + return mem; +} + +inline bool does_size_fit(size_t bytes_per_element, size_t size) +{ + return bytes_per_element * size <= total_global_mem(); +} + +inline std::vector sizes(size_t bytes_per_element) +{ + constexpr size_t all_sizes[] = {1u << 16, 1u << 20, 1u << 24, 1u << 28}; + + std::vector result; + for (size_t size : all_sizes) + { + if (does_size_fit(bytes_per_element, size)) + { + result.push_back(size); + } + } + return result; +} + +} // namespace bench_utils + +PRIMBENCH_REGISTER_TYPE(int8_t, "i8") +PRIMBENCH_REGISTER_TYPE(int16_t, "i16") +PRIMBENCH_REGISTER_TYPE(int32_t, "i32") +PRIMBENCH_REGISTER_TYPE(int64_t, "i64") + +PRIMBENCH_REGISTER_TYPE(uint8_t, "u8") +PRIMBENCH_REGISTER_TYPE(uint16_t, "u16") +PRIMBENCH_REGISTER_TYPE(uint32_t, "u32") +PRIMBENCH_REGISTER_TYPE(uint64_t, "u64") + +#ifndef _MSC_VER +PRIMBENCH_REGISTER_TYPE(int128_t, "i128") +PRIMBENCH_REGISTER_TYPE(uint128_t, "u128") +#endif +PRIMBENCH_REGISTER_TYPE(float, "f32") +PRIMBENCH_REGISTER_TYPE(double, "f64") diff --git a/projects/rocthrust/benchmark/bench_utils/bench_utils.hpp b/projects/rocthrust/benchmark/bench_utils/bench_utils.hpp deleted file mode 100644 index 7899be8cc5c9..000000000000 --- a/projects/rocthrust/benchmark/bench_utils/bench_utils.hpp +++ /dev/null @@ -1,734 +0,0 @@ -// MIT License -// -// Copyright (c) 2024-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 -// in the Software without restriction, including without limitation the rights -// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -// copies of the Software, and to permit persons to whom the Software is -// furnished to do so, subject to the following conditions: -// -// The above copyright notice and this permission notice shall be included in all -// copies or substantial portions of the Software. -// -// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE -// SOFTWARE. - -#ifndef ROCTHRUST_BENCHMARKS_BENCH_UTILS_BENCH_UTILS_HPP_ -#define ROCTHRUST_BENCHMARKS_BENCH_UTILS_BENCH_UTILS_HPP_ - -// Utils -#include - -#include "cmdparser.hpp" -#include "common/types.hpp" -#include "custom_reporter.hpp" -#include "generation_utils.hpp" - -// HIP/CUDA -#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_HIP -# include -#elif THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA -# include -#endif - -// Google Benchmark -#include - -// STL -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -namespace bench_utils -{ -#if (THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HIP) - -# define HIP_CHECK(condition) \ - { \ - hipError_t error = condition; \ - if (error != hipSuccess) \ - { \ - std::cout << "HIP error: " << error << " line: " << __LINE__ << std::endl; \ - exit(error); \ - } \ - } - -/// \brief Timer for measuring time from the device's side -class gpu_timer -{ - hipEvent_t m_start; - hipEvent_t m_stop; - -public: - __forceinline__ gpu_timer() - { - HIP_CHECK(hipEventCreate(&m_start)); - HIP_CHECK(hipEventCreate(&m_stop)); - } - - __forceinline__ ~gpu_timer() - { - HIP_CHECK(hipEventDestroy(m_start)); - HIP_CHECK(hipEventDestroy(m_stop)); - } - - // move-only - gpu_timer(const gpu_timer&) = delete; - gpu_timer(gpu_timer&&) = default; - gpu_timer& operator=(const gpu_timer&) = delete; - gpu_timer& operator=(gpu_timer&&) = default; - - __forceinline__ void start(hipStream_t stream) - { - HIP_CHECK(hipEventRecord(m_start, stream)); - } - - __forceinline__ void stop(hipStream_t stream) - { - HIP_CHECK(hipEventRecord(m_stop, stream)); - } - - [[nodiscard]] __forceinline__ bool ready() const - { - const hipError_t state = hipEventQuery(m_stop); - if (state == hipErrorNotReady) - { - return false; - } - HIP_CHECK(state); - return true; - } - - // In seconds: - [[nodiscard]] __forceinline__ float64_t get_duration() const - { - HIP_CHECK(hipEventSynchronize(m_stop)); - float32_t elapsed_time; - // According to docs, this is in ms with a resolution of ~1 microseconds. - HIP_CHECK(hipEventElapsedTime(&elapsed_time, m_start, m_stop)); - return elapsed_time / 1000.0; - } -}; -#elif (defined(__NVCC__) || defined(_NVHPC_CUDA) \ - || (defined(__CUDA__) && THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_CLANG) \ - || THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_NVRTC) - -# define CUDA_SAFE_CALL_NO_SYNC(call) \ - do \ - { \ - cudaError err = call; \ - if (cudaSuccess != err) \ - { \ - fprintf(stderr, "CUDA error in file '%s' in line %i : %s.\n", __FILE__, __LINE__, cudaGetErrorString(err)); \ - exit(EXIT_FAILURE); \ - } \ - } while (0) - -# define CUDA_SAFE_CALL(call) \ - do \ - { \ - CUDA_SAFE_CALL_NO_SYNC(call); \ - cudaError err = cudaDeviceSynchronize(); \ - if (cudaSuccess != err) \ - { \ - fprintf(stderr, "CUDA error in file '%s' in line %i : %s.\n", __FILE__, __LINE__, cudaGetErrorString(err)); \ - exit(EXIT_FAILURE); \ - } \ - } while (0) - -class gpu_timer -{ - cudaEvent_t start_; - cudaEvent_t stop_; - -public: - __forceinline__ gpu_timer() - { - CUDA_SAFE_CALL(cudaEventCreate(&start_)); - CUDA_SAFE_CALL(cudaEventCreate(&stop_)); - } - - __forceinline__ ~gpu_timer() - { - CUDA_SAFE_CALL(cudaEventDestroy(start_)); - CUDA_SAFE_CALL(cudaEventDestroy(stop_)); - } - - // move-only - gpu_timer(const gpu_timer&) = delete; - gpu_timer(gpu_timer&&) = default; - gpu_timer& operator=(const gpu_timer&) = delete; - gpu_timer& operator=(gpu_timer&&) = default; - - __forceinline__ void start(cudaStream_t stream) - { - CUDA_SAFE_CALL(cudaEventRecord(start_, stream)); - } - - __forceinline__ void stop(cudaStream_t stream) - { - CUDA_SAFE_CALL(cudaEventRecord(m_stop, stream)); - } - - [[nodiscard]] __forceinline__ bool ready() const - { - const cudaError_t state = cudaEventQuery(m_stop); - if (state == cudaErrorNotReady) - { - return false; - } - CUDA_SAFE_CALL(state); - return true; - } - - // In seconds: - [[nodiscard]] __forceinline__ nvbench::float64_t get_duration() const - { - CUDA_SAFE_CALL(cudaEventSynchronize(m_stop)); - float elapsed_time; - // According to docs, this is in ms with a resolution of ~0.5 microseconds. - CUDA_SAFE_CALL(cudaEventElapsedTime(&elapsed_time, m_start, m_stop)); - return elapsed_time / 1000.0; - } -}; - -#endif - -//// \brief Gets the peak global memory bus bandwidth in bytes/sec. -std::size_t get_global_memory_bus_bandwidth(int device_id) -{ - hipDeviceProp_t props; - HIP_CHECK(hipGetDeviceProperties(&props, device_id)); - - // Get the peak clock rate of the global memory bus in Hz. - const std::size_t global_memory_bus_peak_clock_rate = static_cast(props.memoryClockRate) * 1000; /*kHz -> - Hz*/ - // Get width of the global memory bus in bits. - const int get_global_memory_bus_width = props.memoryBusWidth; - - // Multiply by 2 because of DDR, - // CHAR_BIT to convert bus_width to bytes. - return 2 * global_memory_bus_peak_clock_rate * static_cast(get_global_memory_bus_width / CHAR_BIT); -} - -/// \brief Adds device info and properties to the Google benchmark info -inline void add_common_benchmark_info() -{ - hipDeviceProp_t devProp; - int device_id = 0; - HIP_CHECK(hipGetDevice(&device_id)); - HIP_CHECK(hipGetDeviceProperties(&devProp, device_id)); - - auto str = [](const std::string& name, const std::string& val) { - benchmark::AddCustomContext(name, val); - }; - - auto num = [](const std::string& name, const auto& value) { - benchmark::AddCustomContext(name, std::to_string(value)); - }; - - auto dim2 = [num](const std::string& name, const auto* values) { - num(name + "_x", values[0]); - num(name + "_y", values[1]); - }; - - auto dim3 = [num, dim2](const std::string& name, const auto* values) { - dim2(name, values); - num(name + "_z", values[2]); - }; - - str("hdp_name", devProp.name); - num("hdp_total_global_mem", devProp.totalGlobalMem); - num("hdp_shared_mem_per_block", devProp.sharedMemPerBlock); - num("hdp_regs_per_block", devProp.regsPerBlock); - num("hdp_warp_size", devProp.warpSize); - num("hdp_max_threads_per_block", devProp.maxThreadsPerBlock); - dim3("hdp_max_threads_dim", devProp.maxThreadsDim); - dim3("hdp_max_grid_size", devProp.maxGridSize); - num("hdp_clock_rate", devProp.clockRate); - num("hdp_memory_clock_rate", devProp.memoryClockRate); - num("hdp_memory_bus_width", devProp.memoryBusWidth); - num("hdp_peak_global_mem_bus_bandwidth", get_global_memory_bus_bandwidth(device_id)); - num("hdp_total_const_mem", devProp.totalConstMem); - num("hdp_major", devProp.major); - num("hdp_minor", devProp.minor); - num("hdp_multi_processor_count", devProp.multiProcessorCount); - num("hdp_l2_cache_size", devProp.l2CacheSize); - num("hdp_max_threads_per_multiprocessor", devProp.maxThreadsPerMultiProcessor); - num("hdp_compute_mode", devProp.computeMode); - num("hdp_clock_instruction_rate", devProp.clockInstructionRate); - num("hdp_concurrent_kernels", devProp.concurrentKernels); - num("hdp_pci_domain_id", devProp.pciDomainID); - num("hdp_pci_bus_id", devProp.pciBusID); - num("hdp_pci_device_id", devProp.pciDeviceID); - num("hdp_max_shared_memory_per_multi_processor", devProp.maxSharedMemoryPerMultiProcessor); - num("hdp_is_multi_gpu_board", devProp.isMultiGpuBoard); - num("hdp_can_map_host_memory", devProp.canMapHostMemory); - str("hdp_gcn_arch_name", devProp.gcnArchName); - num("hdp_integrated", devProp.integrated); - num("hdp_cooperative_launch", devProp.cooperativeLaunch); - num("hdp_cooperative_multi_device_launch", devProp.cooperativeMultiDeviceLaunch); - num("hdp_max_texture_1d_linear", devProp.maxTexture1DLinear); - num("hdp_max_texture_1d", devProp.maxTexture1D); - dim2("hdp_max_texture_2d", devProp.maxTexture2D); - dim3("hdp_max_texture_3d", devProp.maxTexture3D); - num("hdp_mem_pitch", devProp.memPitch); - num("hdp_texture_alignment", devProp.textureAlignment); - num("hdp_texture_pitch_alignment", devProp.texturePitchAlignment); - num("hdp_kernel_exec_timeout_enabled", devProp.kernelExecTimeoutEnabled); - num("hdp_ecc_enabled", devProp.ECCEnabled); - num("hdp_tcc_driver", devProp.tccDriver); - num("hdp_cooperative_multi_device_unmatched_func", devProp.cooperativeMultiDeviceUnmatchedFunc); - num("hdp_cooperative_multi_device_unmatched_grid_dim", devProp.cooperativeMultiDeviceUnmatchedGridDim); - num("hdp_cooperative_multi_device_unmatched_block_dim", devProp.cooperativeMultiDeviceUnmatchedBlockDim); - num("hdp_cooperative_multi_device_unmatched_shared_mem", devProp.cooperativeMultiDeviceUnmatchedSharedMem); - num("hdp_is_large_bar", devProp.isLargeBar); - num("hdp_asic_revision", devProp.asicRevision); - num("hdp_managed_memory", devProp.managedMemory); - num("hdp_direct_managed_mem_access_from_host", devProp.directManagedMemAccessFromHost); - num("hdp_concurrent_managed_access", devProp.concurrentManagedAccess); - num("hdp_pageable_memory_access", devProp.pageableMemoryAccess); - num("hdp_pageable_memory_access_uses_host_page_tables", devProp.pageableMemoryAccessUsesHostPageTables); - - const auto arch = devProp.arch; - num("hdp_arch_has_global_int32_atomics", arch.hasGlobalInt32Atomics); - num("hdp_arch_has_global_float_atomic_exch", arch.hasGlobalFloatAtomicExch); - num("hdp_arch_has_shared_int32_atomics", arch.hasSharedInt32Atomics); - num("hdp_arch_has_shared_float_atomic_exch", arch.hasSharedFloatAtomicExch); - num("hdp_arch_has_float_atomic_add", arch.hasFloatAtomicAdd); - num("hdp_arch_has_global_int64_atomics", arch.hasGlobalInt64Atomics); - num("hdp_arch_has_shared_int64_atomics", arch.hasSharedInt64Atomics); - num("hdp_arch_has_doubles", arch.hasDoubles); - num("hdp_arch_has_warp_vote", arch.hasWarpVote); - num("hdp_arch_has_warp_ballot", arch.hasWarpBallot); - num("hdp_arch_has_warp_shuffle", arch.hasWarpShuffle); - num("hdp_arch_has_funnel_shift", arch.hasFunnelShift); - num("hdp_arch_has_thread_fence_system", arch.hasThreadFenceSystem); - num("hdp_arch_has_sync_threads_ext", arch.hasSyncThreadsExt); - num("hdp_arch_has_surface_funcs", arch.hasSurfaceFuncs); - num("hdp_arch_has_3d_grid", arch.has3dGrid); - num("hdp_arch_has_dynamic_parallelism", arch.hasDynamicParallelism); -} - -// Binary operators -struct less_t -{ - template - __host__ __device__ bool operator()(const T& lhs, const T& rhs) const - { - return lhs < rhs; - } -}; - -struct max_t -{ - template - __host__ __device__ T operator()(const T& lhs, const T& rhs) - { - less_t less{}; - return less(lhs, rhs) ? rhs : lhs; - } -}; - -struct bench_naming -{ -public: - enum format - { - json, - human, - txt - }; - static format& get_format() - { - static format storage = human; - return storage; - } - static void set_format(const std::string& argument) - { - format result = human; - if (argument == "json") - { - result = json; - } - else if (argument == "txt") - { - result = txt; - } - get_format() = result; - } - -private: - static std::string matches_as_json(std::sregex_iterator& matches) - { - std::stringstream result; - int brackets_count = 1; - result << "{"; - bool insert_comma = false; - for (std::sregex_iterator i = matches; i != std::sregex_iterator(); ++i) - { - std::smatch m = *i; - if (insert_comma) - { - result << ","; - } - else - { - insert_comma = true; - } - result << "\"" << m[1].str() << "\":"; - if (m[2].length() > 0) - { - if (m[2].str().find_first_not_of("0123456789") == std::string::npos) - { - result << m[2].str(); - } - else - { - result << "\"" << m[2].str() << "\""; - } - if (m[3].length() > 0 && brackets_count > 0) - { - int n = std::min(brackets_count, static_cast(m[3].length())); - brackets_count -= n; - for (int c = 0; c < n; c++) - { - result << "}"; - } - } - } - else - { - brackets_count++; - result << "{"; - insert_comma = false; - } - } - while (brackets_count > 0) - { - brackets_count--; - result << "}"; - } - return result.str(); - } - - static std::string matches_as_human(std::sregex_iterator& matches) - { - std::stringstream result; - int brackets_count = 0; - bool insert_comma = false; - for (std::sregex_iterator i = matches; i != std::sregex_iterator(); ++i) - { - std::smatch m = *i; - if (insert_comma) - { - result << ","; - } - else - { - insert_comma = true; - } - if (m[2].length() > 0) - { - result << m[2].str(); - if (m[3].length() > 0 && brackets_count > 0) - { - int n = std::min(brackets_count, static_cast(m[3].length())); - brackets_count -= n; - for (int c = 0; c < n; c++) - { - result << ">"; - } - } - } - else - { - brackets_count++; - result << "<"; - insert_comma = false; - } - } - while (brackets_count > 0) - { - brackets_count--; - result << ">"; - } - return result.str(); - } - -public: - static std::string format_name(std::string string) - { - format format = get_format(); - std::regex r("([A-z0-9_]*):([A-z_:\\(\\)\\.<>\\s0-9\" ]*)"); - // First we perform some checks - bool checks[5] = {false}; - for (std::sregex_iterator i = std::sregex_iterator(string.begin(), string.end(), r); i != std::sregex_iterator(); - ++i) - { - std::smatch m = *i; - if (m[1].str() == "algo") - { - checks[0] = true; - } - else if (m[1].str() == "subalgo") - { - checks[1] = true; - } - else if (m[1].str() == "input_type" || m[1].str() == "key_type" || m[1].str() == "value_type") - { - checks[2] = true; - } - else if (m[1].str() == "elements") - { - checks[3] = true; - } - } - std::string string_substitute = std::regex_replace(string, r, ""); - checks[4] = string_substitute.find_first_not_of(" ,{}") == std::string::npos; - for (bool check_name_format : checks) - { - if (!check_name_format) - { - std::cout << "string_substitute = " << string_substitute << std::endl; - std::cout - << "Benchmark name \"" << string - << "\" not in the correct format (e.g. " - "{algo:reduce,subalgo:by_key} )" - << std::endl; - exit(1); - } - } - - // Now we generate the desired format - std::sregex_iterator matches = std::sregex_iterator(string.begin(), string.end(), r); - - switch (format) - { - case format::json: - return matches_as_json(matches); - case format::human: - return matches_as_human(matches); - case format::txt: - return string; - } - return string; - } -}; - -namespace detail -{ -void do_not_optimize(const void* ptr) -{ - (void) ptr; -} -} // namespace detail - -template -void do_not_optimize(const T& val) -{ - detail::do_not_optimize(&val); -} - -auto StatisticsSum = [](const std::vector& v) { - return std::accumulate(v.begin(), v.end(), 0.0); -}; - -double StatisticsMean(const std::vector& v) -{ - if (v.empty()) - { - return 0.0; - } - return StatisticsSum(v) * (1.0 / static_cast(v.size())); -} - -double StatisticsMedian(const std::vector& v) -{ - if (v.size() < 3) - { - return StatisticsMean(v); - } - std::vector copy(v); - - auto center = copy.begin() + v.size() / 2; - std::nth_element(copy.begin(), center, copy.end()); - - // Did we have an odd number of samples? If yes, then center is the median. - // If not, then we are looking for the average between center and the value - // before. Instead of resorting, we just look for the max value before it, - // which is not necessarily the element immediately preceding `center` Since - // `copy` is only partially sorted by `nth_element`. - if (v.size() % 2 == 1) - { - return *center; - } - auto center2 = std::max_element(copy.begin(), center); - return (*center + *center2) / 2.0; -} - -// Return the sum of the squares of this sample set -auto SumSquares = [](const std::vector& v) { - return std::inner_product(v.begin(), v.end(), v.begin(), 0.0); -}; - -auto Sqr = [](const double dat) { - return dat * dat; -}; -auto Sqrt = [](const double dat) { - // Avoid NaN due to imprecision in the calculations - if (dat < 0.0) - { - return 0.0; - } - return std::sqrt(dat); -}; - -double StatisticsStdDev(const std::vector& v) -{ - const auto mean = StatisticsMean(v); - if (v.empty()) - { - return mean; - } - - // Sample standard deviation is undefined for n = 1 - if (v.size() == 1) - { - return 0.0; - } - - const double avg_squares = SumSquares(v) * (1.0 / static_cast(v.size())); - return Sqrt(static_cast(v.size()) / (static_cast(v.size()) - 1.0) * (avg_squares - Sqr(mean))); -} - -double StatisticsCV(const std::vector& v) -{ - if (v.size() < 2) - { - return 0.0; - } - - const auto stddev = StatisticsStdDev(v); - const auto mean = StatisticsMean(v); - - if (std::fpclassify(mean) == FP_ZERO) - { - return 0.0; - } - - return stddev / mean; -} - -inline const char* get_seed_message() -{ - return "seed for input generation, either an unsigned integer value for determinisic results " - "or 'random' for different inputs for each repetition"; -} - -struct caching_allocator_t -{ - using value_type = char; - - caching_allocator_t() = default; - ~caching_allocator_t() - { - free_all(); - } - - char* allocate(std::ptrdiff_t num_bytes) - { - value_type* result{}; - auto free_block = free_blocks.find(num_bytes); - if (free_block != free_blocks.end()) - { - result = free_block->second; - free_blocks.erase(free_block); - } - else - { - HIP_CHECK(hipMalloc(&result, num_bytes)); - } - - allocated_blocks.emplace(result, num_bytes); - return result; - } - - void deallocate(value_type* ptr, size_t) - { - auto iter = allocated_blocks.find(ptr); - if (iter == allocated_blocks.end()) - { - throw std::runtime_error("Memory was not allocated by this allocator"); - } - - std::ptrdiff_t num_bytes = iter->second; - allocated_blocks.erase(iter); - free_blocks.emplace(num_bytes, ptr); - } - -private: - using FreeBlocksType = std::multimap; - using AllocatedBlocksType = std::map; - - FreeBlocksType free_blocks; - AllocatedBlocksType allocated_blocks; - - void free_all() - { - for (auto free_block : free_blocks) - { - HIP_CHECK(hipFree(free_block.second)); - } - - for (auto allocated_block : allocated_blocks) - { - HIP_CHECK(hipFree(allocated_block.first)); - } - } -}; - -inline std::string format_pow2(size_t n) -{ - unsigned int k = 0; - while (!(n & 1)) - { - k++; - n >>= 1; - } - return "1 << " + std::to_string(k); -} - -struct sys_info -{ - hipDeviceProp_t devProp; - sys_info() - { - int device_id = 0; - HIP_CHECK(hipGetDevice(&device_id)); - HIP_CHECK(hipGetDeviceProperties(&devProp, device_id)); - } -}; - -inline sys_info system; -inline constexpr size_t sizes[] = {1u << 16, 1u << 20, 1u << 24, 1u << 28}; - -} // namespace bench_utils - -#endif // ROCTHRUST_BENCHMARKS_BENCH_UTILS_BENCH_UTILS_HPP_ diff --git a/projects/rocthrust/benchmark/bench_utils/cmdparser.hpp b/projects/rocthrust/benchmark/bench_utils/cmdparser.hpp deleted file mode 100644 index da7724140eab..000000000000 --- a/projects/rocthrust/benchmark/bench_utils/cmdparser.hpp +++ /dev/null @@ -1,657 +0,0 @@ -// The MIT License (MIT) -// -// Copyright (c) 2015 - 2016 Florian Rappl -// Modifications Copyright (c) 2024-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 -// in the Software without restriction, including without limitation the rights -// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -// copies of the Software, and to permit persons to whom the Software is -// furnished to do so, subject to the following conditions: -// -// The above copyright notice and this permission notice shall be included in all -// copies or substantial portions of the Software. -// -// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE -// SOFTWARE. - -/* - This file is part of the C++ CmdParser utility. - Copyright (c) 2015 - 2016 Florian Rappl -*/ - -#pragma once -#include -#include -#include -#include -#include -#include - -namespace cli -{ -struct CallbackArgs -{ - const std::vector& arguments; - std::ostream& output; - std::ostream& error; -}; -class Parser -{ -private: - class CmdBase - { - public: - explicit CmdBase(const std::string& name, - const std::string& alternative, - const std::string& description, - bool required, - bool dominant, - bool variadic) - : name(name) - , command(name.size() > 0 ? "-" + name : "") - , alternative(alternative.size() > 0 ? "--" + alternative : "") - , description(description) - , required(required) - , handled(false) - , arguments({}) - , dominant(dominant) - , variadic(variadic) - {} - - virtual ~CmdBase() {} - - std::string name; - std::string command; - std::string alternative; - std::string description; - bool required; - bool handled; - std::vector arguments; - bool const dominant; - bool const variadic; - - virtual std::string print_value() const = 0; - virtual bool parse(std::ostream& output, std::ostream& error) = 0; - - bool is(const std::string& given) const - { - return given == command || given == alternative; - } - }; - - template - struct ArgumentCountChecker - { - static constexpr bool Variadic = false; - }; - - template - struct ArgumentCountChecker> - { - static constexpr bool Variadic = true; - }; - - template - class CmdFunction final : public CmdBase - { - public: - explicit CmdFunction(const std::string& name, - const std::string& alternative, - const std::string& description, - bool required, - bool dominant) - : CmdBase(name, alternative, description, required, dominant, ArgumentCountChecker::Variadic) - {} - - virtual bool parse(std::ostream& output, std::ostream& error) override - { - try - { - CallbackArgs args{arguments, output, error}; - value = callback(args); - return true; - } - catch (...) - { - return false; - } - } - - virtual std::string print_value() const override - { - return ""; - } - - std::function callback; - T value; - }; - - template - class CmdArgument final : public CmdBase - { - public: - explicit CmdArgument(const std::string& name, - const std::string& alternative, - const std::string& description, - bool required, - bool dominant) - : CmdBase(name, alternative, description, required, dominant, ArgumentCountChecker::Variadic) - , value(T()) - {} - - virtual bool parse(std::ostream&, std::ostream&) override - { - try - { - value = Parser::parse(arguments, value); - return true; - } - catch (...) - { - return false; - } - } - - virtual std::string print_value() const override - { - return stringify(value); - } - - T value; - }; - - static int parse(const std::vector& elements, const int&) - { - if (elements.size() != 1) - { - throw std::bad_cast(); - } - - return std::stoi(elements[0]); - } - - static bool parse(const std::vector& elements, const bool& defval) - { - if (elements.size() != 0) - { - throw std::runtime_error("A boolean command line parameter cannot have any arguments."); - } - - return !defval; - } - - static double parse(const std::vector& elements, const double&) - { - if (elements.size() != 1) - { - throw std::bad_cast(); - } - - return std::stod(elements[0]); - } - - static float parse(const std::vector& elements, const float&) - { - if (elements.size() != 1) - { - throw std::bad_cast(); - } - - return std::stof(elements[0]); - } - - static long double parse(const std::vector& elements, const long double&) - { - if (elements.size() != 1) - { - throw std::bad_cast(); - } - - return std::stold(elements[0]); - } - - static unsigned int parse(const std::vector& elements, const unsigned int&) - { - if (elements.size() != 1) - { - throw std::bad_cast(); - } - - return static_cast(std::stoul(elements[0])); - } - - static unsigned long parse(const std::vector& elements, const unsigned long&) - { - if (elements.size() != 1) - { - throw std::bad_cast(); - } - - return std::stoul(elements[0]); - } - - static unsigned long long parse(const std::vector& elements, const unsigned long long&) - { - if (elements.size() != 1) - { - throw std::bad_cast(); - } - - return std::stoull(elements[0]); - } - - static long parse(const std::vector& elements, const long&) - { - if (elements.size() != 1) - { - throw std::bad_cast(); - } - - return std::stol(elements[0]); - } - - static std::string parse(const std::vector& elements, const std::string&) - { - if (elements.size() != 1) - { - throw std::bad_cast(); - } - - return elements[0]; - } - - template - static std::vector parse(const std::vector& elements, const std::vector&) - { - const T defval = T(); - std::vector values{}; - std::vector buffer(1); - - for (const auto& element : elements) - { - buffer[0] = element; - values.push_back(parse(buffer, defval)); - } - - return values; - } - - template - static std::string stringify(const T& value) - { - return std::to_string(value); - } - - template - static std::string stringify(const std::vector& values) - { - std::stringstream ss{}; - ss << "[ "; - - for (const auto& value : values) - { - ss << stringify(value) << " "; - } - - ss << "]"; - return ss.str(); - } - - static std::string stringify(const std::string& str) - { - return str; - } - -public: - explicit Parser(int argc, const char** argv) - : _appname(argv[0]) - { - for (int i = 1; i < argc; ++i) - { - _arguments.push_back(argv[i]); - } - enable_help(); - } - - explicit Parser(int argc, char** argv) - : _appname(argv[0]) - { - for (int i = 1; i < argc; ++i) - { - _arguments.push_back(argv[i]); - } - enable_help(); - } - - ~Parser() - { - for (int i = 0, n = _commands.size(); i < n; ++i) - { - delete _commands[i]; - } - } - - bool has_help() const - { - for (const auto command : _commands) - { - if (command->name == "h" && command->alternative == "--help") - { - return true; - } - } - - return false; - } - - void enable_help() - { - set_callback( - "h", - "help", - std::function([this](CallbackArgs& args) { - args.output << this->usage(); - /*exit(0);*/ - return false; - }), - "", - true); - } - - void disable_help() - { - for (auto command = _commands.begin(); command != _commands.end(); ++command) - { - if ((*command)->name == "h" && (*command)->alternative == "--help") - { - _commands.erase(command); - break; - } - } - } - - template - void set_default(bool is_required, const std::string& description = "") - { - auto command = new CmdArgument{"", "", description, is_required, false}; - _commands.push_back(command); - } - - template - void set_required( - const std::string& name, const std::string& alternative, const std::string& description = "", bool dominant = false) - { - auto command = new CmdArgument{name, alternative, description, true, dominant}; - _commands.push_back(command); - } - - template - void set_optional(const std::string& name, - const std::string& alternative, - const T& defaultValue, - const std::string& description = "", - bool dominant = false) - { - auto command = new CmdArgument{name, alternative, description, false, dominant}; - command->value = defaultValue; - _commands.push_back(command); - } - - template - void set_callback(const std::string& name, - const std::string& alternative, - std::function callback, - const std::string& description = "", - bool dominant = false) - { - auto command = new CmdFunction{name, alternative, description, false, dominant}; - command->callback = callback; - _commands.push_back(command); - } - - inline void run_and_exit_if_error() - { - if (run() == false) - { - exit(1); - } - } - - inline bool run() - { - return run(std::cout, std::cerr); - } - - inline bool run(std::ostream& output) - { - return run(output, std::cerr); - } - - bool run(std::ostream& output, std::ostream& error) - { - if (_arguments.size() > 0) - { - auto current = find_default(); - - for (int i = 0, n = _arguments.size(); i < n; ++i) - { - auto isarg = _arguments[i].size() > 0 && _arguments[i][0] == '-'; - auto associated = isarg ? find(_arguments[i]) : nullptr; - - if (associated != nullptr) - { - current = associated; - associated->handled = true; - } - else if (current == nullptr) - { - current = find(_arguments[i]); - // Code was commented out so cmdparser can ignore unknown options - // error << no_default(); - // return false; - } - else - { - current->arguments.push_back(_arguments[i]); - current->handled = true; - if (!current->variadic) - { - // If the current command is not variadic, then no more arguments - // should be added to it. In this case, switch back to the default - // command. - current = find_default(); - } - } - } - } - - // First, parse dominant arguments since they succeed even if required - // arguments are missing. - for (auto command : _commands) - { - if (command->handled && command->dominant && !command->parse(output, error)) - { - error << howto_use(command); - return false; - } - } - - // Next, check for any missing arguments. - for (auto command : _commands) - { - if (command->required && !command->handled) - { - error << howto_required(command); - return false; - } - } - - // Finally, parse all remaining arguments. - for (auto command : _commands) - { - if (command->handled && !command->dominant && !command->parse(output, error)) - { - error << howto_use(command); - return false; - } - } - - return true; - } - - template - T get(const std::string& name) const - { - for (const auto& command : _commands) - { - if (command->name == name) - { - auto cmd = dynamic_cast*>(command); - - if (cmd == nullptr) - { - throw std::runtime_error("Invalid usage of the parameter " + name + " detected."); - } - - return cmd->value; - } - } - - throw std::runtime_error("The parameter " + name + " could not be found."); - } - - template - T get_if(const std::string& name, std::function callback) const - { - auto value = get(name); - return callback(value); - } - - int requirements() const - { - int count = 0; - - for (const auto& command : _commands) - { - if (command->required) - { - ++count; - } - } - - return count; - } - - int commands() const - { - return static_cast(_commands.size()); - } - - inline const std::string& app_name() const - { - return _appname; - } - -protected: - CmdBase* find(const std::string& name) - { - for (auto command : _commands) - { - if (command->is(name)) - { - return command; - } - } - - return nullptr; - } - - CmdBase* find_default() - { - for (auto command : _commands) - { - if (command->name == "") - { - return command; - } - } - - return nullptr; - } - - std::string usage() const - { - std::stringstream ss{}; - ss << "Available parameters:\n\n"; - - for (const auto& command : _commands) - { - ss << " " << command->command << "\t" << command->alternative; - - if (command->required == true) - { - ss << "\t(required)"; - } - - ss << "\n " << command->description; - - if (command->required == false) - { - ss << "\n " - << "This parameter is optional. The default value is '" + command->print_value() << "'."; - } - - ss << "\n\n"; - } - - return ss.str(); - } - - void print_help(std::stringstream& ss) const - { - if (has_help()) - { - ss << "For more help use --help or -h.\n"; - } - } - - std::string howto_required(CmdBase* command) const - { - std::stringstream ss{}; - ss << "The parameter " << command->name << " is required.\n"; - ss << command->description << '\n'; - print_help(ss); - return ss.str(); - } - - std::string howto_use(CmdBase* command) const - { - std::stringstream ss{}; - ss << "The parameter " << command->name << " has invalid arguments.\n"; - ss << command->description << '\n'; - print_help(ss); - return ss.str(); - } - - std::string no_default() const - { - std::stringstream ss{}; - ss << "No default parameter has been specified.\n"; - ss << "The given argument must be used with a parameter.\n"; - print_help(ss); - return ss.str(); - } - -private: - const std::string _appname; - std::vector _arguments; - std::vector _commands; -}; -} // namespace cli diff --git a/projects/rocthrust/benchmark/bench_utils/common/types.hpp b/projects/rocthrust/benchmark/bench_utils/common/types.hpp deleted file mode 100644 index 3eaa2c925b85..000000000000 --- a/projects/rocthrust/benchmark/bench_utils/common/types.hpp +++ /dev/null @@ -1,97 +0,0 @@ -// MIT License -// -// Copyright (c) 2024-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 -// in the Software without restriction, including without limitation the rights -// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -// copies of the Software, and to permit persons to whom the Software is -// furnished to do so, subject to the following conditions: -// -// The above copyright notice and this permission notice shall be included in all -// copies or substantial portions of the Software. -// -// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE -// SOFTWARE. - -#ifndef ROCTHRUST_BENCHMARKS_BENCH_UTILS_TYPES_HPP_ -#define ROCTHRUST_BENCHMARKS_BENCH_UTILS_TYPES_HPP_ - -#include - -// Types used in the benchmarks -#if defined(_MSC_VER) -# define THRUST_BENCHMARKS_HAVE_INT128_SUPPORT 0 -#else -# define THRUST_BENCHMARKS_HAVE_INT128_SUPPORT 1 -#endif - -namespace bench_utils -{ -class large_data -{ -public: - __host__ __device__ large_data() - { - data[0] = 0; - } - __host__ __device__ large_data(large_data const& val) - { - data[0] = val.data[0]; - } - __host__ __device__ large_data(int n) - { - data[0] = static_cast(n); - } - large_data& __host__ __device__ operator=(large_data const& val) - { - data[0] = val.data[0]; - return *this; - } - bool __host__ __device__ operator==(large_data const& val) const - { - return data[0] == val.data[0]; - } - large_data& __host__ __device__ operator++() - { - ++data[0]; - return *this; - } - __host__ __device__ operator int() const - { - return static_cast(data[0]); - } - - int8_t data[512]; -}; - -template -bool __host__ __device__ operator==(T const& lhs, large_data const& rhs) -{ - return static_cast(lhs).data[0] == rhs.data[0]; -} - -}; // namespace bench_utils - -using int8_t = std::int8_t; -using int16_t = std::int16_t; -using int32_t = std::int32_t; -using int64_t = std::int64_t; -using uint8_t = std::uint8_t; -using uint16_t = std::uint16_t; -using uint32_t = std::uint32_t; -using uint64_t = std::uint64_t; -#if THRUST_BENCHMARKS_HAVE_INT128_SUPPORT -using int128_t = __int128_t; -using uint128_t = __uint128_t; -#endif -using float32_t = float; -using float64_t = double; - -#endif // ROCTHRUST_BENCHMARKS_BENCH_UTILS_TYPES_HPP_ diff --git a/projects/rocthrust/benchmark/bench_utils/custom_reporter.hpp b/projects/rocthrust/benchmark/bench_utils/custom_reporter.hpp deleted file mode 100644 index 986c8d57ee1e..000000000000 --- a/projects/rocthrust/benchmark/bench_utils/custom_reporter.hpp +++ /dev/null @@ -1,826 +0,0 @@ -// MIT License -// -// Copyright (c) 2024-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 -// in the Software without restriction, including without limitation the rights -// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -// copies of the Software, and to permit persons to whom the Software is -// furnished to do so, subject to the following conditions: -// -// The above copyright notice and this permission notice shall be included in all -// copies or substantial portions of the Software. -// -// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE -// SOFTWARE. - -#ifndef ROCTHRUST_BENCHMARKS_BENCH_UTILS_CUSTOM_REPORTER_HPP_ -#define ROCTHRUST_BENCHMARKS_BENCH_UTILS_CUSTOM_REPORTER_HPP_ - -// Utils -#include "common/types.hpp" - -// Google Benchmark -#include - -// STL -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -namespace bench_utils -{ - -template -bool IsType(const SrcType* src) -{ - // Check if the src can be casted to the DstType - return dynamic_cast(src) != nullptr; -} - -static std::string FormatString(const char* msg, va_list args) -{ - // we might need a second shot at this, so pre-emptivly make a copy - va_list args_cp; - va_copy(args_cp, args); - - std::size_t size = 256; - char local_buff[256]; - auto ret = vsnprintf(local_buff, size, msg, args_cp); - - if (ret <= 0) - { - return {}; - } - else if (static_cast(ret) < size) - { - return local_buff; - } - - // we did not provide a long enough buffer on our first attempt. - size = static_cast(ret) + 1; // + 1 for the null byte - std::unique_ptr buff(new char[size]); - ret = vsnprintf(buff.get(), size, msg, args); - return buff.get(); -} - -static std::string FormatString(const char* msg, ...) -{ - va_list args; - va_start(args, msg); - auto tmp = FormatString(msg, args); - va_end(args); - return tmp; -} - -std::string get_complexity(const benchmark::BigO& complexity) -{ - switch (complexity) - { - case benchmark::oN: - return "N"; - case benchmark::oNSquared: - return "N^2"; - case benchmark::oNCubed: - return "N^3"; - case benchmark::oLogN: - return "lgN"; - case benchmark::oNLogN: - return "NlgN"; - case benchmark::o1: - return "(1)"; - default: - return "f(N)"; - } -} - -template -double calculate_bw_utils(const Run& result) -{ - // Calculates bandwith utilization in % - std::map* global_context = benchmark::internal::GetGlobalContext(); - if (global_context != nullptr) - { - for (const auto& keyval : *global_context) - { - if (keyval.first == "hdp_peak_global_mem_bus_bandwidth") - { - const double global_mem_bw = result.counters.at("bytes_per_second").value; - const double peak_global_bw = std::stod(keyval.second); - const double bw_util = 100. * global_mem_bw / peak_global_bw; - return bw_util; - } - } - } - return -1; -} - -/// \brief Custom Google Benchmark reporter for formatting the benchmarks' report matching Thrust's. -/// -/// This reporter is a ConsoleReporter that outputs: -/// - GPU Time: measured with events and registered as manual time. -/// - CPU Time: measured by Google Benchmark. -/// - Iterations: the number of kernel executions done on each benchmark run. -/// - GlobalMem BW: number of GB read and written to global memory per second of execution. -/// - BW util: percentage of the theoretical peak global memory bandwith utilised by the benchmark. -/// - Elements/s: number of G elements (G = 10**6) processed by second of execution. -/// - Labels: for extra labels. For instance, labels are used by the benchmarks generating random -/// input values to report the seed used. -/// -/// Additionally, when the number of \p repetitions is greater than one, each benchmark run is -/// repeated \p repetitions times to measure the stability of results. In this case, the mean, -/// median, standard deviation (stddev) and coefficient of variation (cv) of the above-described -/// metrics are also reported after all the \p repetitions have ben run. -class CustomConsoleReporter : public benchmark::ConsoleReporter -{ -private: - enum LogColor - { - COLOR_DEFAULT, - COLOR_RED, - COLOR_GREEN, - COLOR_YELLOW, - COLOR_BLUE, - COLOR_CYAN, - COLOR_WHITE - }; - - std::string get_log_color(LogColor&& color) - { - std::string s("unknown"); - switch (color) - { - case COLOR_RED: - return "\033[31m"; - case COLOR_GREEN: - return "\033[32m"; - case COLOR_YELLOW: - return "\033[33m"; - case COLOR_BLUE: - return "\033[34m"; - case COLOR_CYAN: - return "\033[36m"; - case COLOR_WHITE: - return "\033[37m"; - default: // COLOR_DEFAULT - return "\033[0m"; - } - } - - void PrintColoredString(std::ostream& os, std::string color, std::string str, ...) - { - os << color; - va_list args; - va_start(args, str); - os << FormatString(str.data(), args); - va_end(args); - } - - static std::string FormatTime(double time) - { - // Assuming the time is at max 9.9999e+99 and we have 10 digits for the - // number, we get 10-1(.)-1(e)-1(sign)-2(exponent) = 5 digits to print. - if (time > 9999999999 /*max 10 digit number*/) - { - return FormatString("%1.4e", time); - } - return FormatString("%10.3f", time); - } - -public: - void PrintHeader(const Run& /*run*/) - { - // Assume run.counters (with elements processed and global mem reads/writes) - // will not be empty - std::string str = FormatString( - "%-*s %13s %15s %12s %16s %12s %14s %14s", - static_cast(name_field_width_), - "Benchmark", - "GPU Time", - "CPU Time", - "Iterations", - "GlobalMem BW", - "BW util", - "GPU Noise", - "Elements/s"); - std::string line = std::string(str.length(), '-'); - GetOutputStream() << line << "\n" << str << "\n" << line << "\n"; - } - - void PrintRunData(const Run& result) - { - // Report benchmark name - auto& sout = GetOutputStream(); - auto color = get_log_color((result.report_big_o || result.report_rms) ? COLOR_BLUE : COLOR_GREEN); - - PrintColoredString(sout, color, "%-*s ", name_field_width_, result.benchmark_name().c_str()); - - if (benchmark::internal::SkippedWithError == result.skipped) - { - PrintColoredString(sout, get_log_color(COLOR_RED), "ERROR OCCURRED: \'%s\'", result.skip_message.c_str()); - PrintColoredString(sout, get_log_color(COLOR_DEFAULT), "\n"); - return; - } - else if (benchmark::internal::SkippedWithMessage == result.skipped) - { - PrintColoredString(sout, get_log_color(COLOR_WHITE), "SKIPPED: \'%s\'", result.skip_message.c_str()); - PrintColoredString(sout, get_log_color(COLOR_DEFAULT), "\n"); - return; - } - - // Report GPU and CPU time - const double gpu_time = result.GetAdjustedRealTime(); - const double cpu_time = result.GetAdjustedCPUTime(); - const std::string gpu_time_str = FormatTime(gpu_time); - const std::string cpu_time_str = FormatTime(cpu_time); - - if (result.report_big_o) - { - std::string big_o = get_complexity(result.complexity); - PrintColoredString( - sout, get_log_color(COLOR_YELLOW), "%10.3f %-4s %10.3f %-4s ", gpu_time, big_o.c_str(), cpu_time, big_o.c_str()); - } - else if (result.report_rms) - { - PrintColoredString( - sout, get_log_color(COLOR_YELLOW), "%10.3f %-4s %10.3f %-4s ", gpu_time * 100, "%", cpu_time * 100, "%"); - } - else if (result.run_type != Run::RT_Aggregate || result.aggregate_unit == benchmark::StatisticUnit::kTime) - { - const char* timeLabel = GetTimeUnitString(result.time_unit); - PrintColoredString( - sout, - get_log_color(COLOR_YELLOW), - "%s %-4s %s %-4s ", - gpu_time_str.c_str(), - timeLabel, - cpu_time_str.c_str(), - timeLabel); - } - else - { - assert(result.aggregate_unit == benchmark::StatisticUnit::kPercentage); - PrintColoredString( - sout, - get_log_color(COLOR_YELLOW), - "%10.3f %-4s %10.3f %-4s ", - (100. * result.real_accumulated_time), - "%", - (100. * result.cpu_accumulated_time), - "%"); - } - - // Report iterations - if (!result.report_big_o && !result.report_rms) - { - PrintColoredString(sout, get_log_color(COLOR_CYAN), "%10lld", result.iterations); - } - - // Report counters - std::string s; - const char* unit = ""; - for (auto& c : result.counters) - { - std::size_t cNameLen = std::max(std::string::size_type(10), c.first.length()); - if (c.first == "items_per_second") - { - // Print elements processed in G/s - if (result.run_type == Run::RT_Aggregate && result.aggregate_unit == benchmark::StatisticUnit::kPercentage) - { - s = FormatString("%.3f", 100. * c.second.value); - unit = "%"; - } - else - { - s = FormatString("%.3f", c.second.value / 1e9); - unit = "G"; - } - } - else if (c.first == "bytes_per_second") - { - // Print GlobalMem BW in GB/s - if (result.run_type == Run::RT_Aggregate && result.aggregate_unit == benchmark::StatisticUnit::kPercentage) - { - s = FormatString("%.3f", 100. * c.second.value); - unit = "%"; - } - else - { - s = FormatString("%.3f", c.second.value / 1e9); - if (c.second.flags & benchmark::Counter::kIsRate) - { - unit = "GB/s"; - } - } - PrintColoredString(sout, get_log_color(COLOR_DEFAULT), " %*s%s", cNameLen - strlen(unit), s.c_str(), unit); - - // Print BW util - s = FormatString("%.2f", c.second.value); - const double bw_util = calculate_bw_utils(result); - if (bw_util >= 0) - { - s = FormatString("%.2f", bw_util); - unit = "%"; - cNameLen = std::max(std::string::size_type(12), s.length()); - } - } - else if (c.first == "gpu_noise") - { - if (result.run_type == Run::RT_Aggregate && result.aggregate_unit == benchmark::StatisticUnit::kPercentage) - { - s = FormatString("%.2f", 100. * c.second.value); - } - else - { - s = FormatString("%.2f", c.second.value); - } - unit = "%"; - } - else - { - if (result.run_type == Run::RT_Aggregate && result.aggregate_unit == benchmark::StatisticUnit::kPercentage) - { - s = FormatString("%.2f", 100. * c.second.value); - unit = "%"; - } - else - { - s = FormatString("%.2f", c.second.value / 1000.); - if (c.second.flags & benchmark::Counter::kIsRate) - { - unit = (c.second.flags & benchmark::Counter::kInvert) ? "s" : "/s"; - } - } - } - - PrintColoredString(sout, get_log_color(COLOR_DEFAULT), " %*s%s", cNameLen - strlen(unit), s.c_str(), unit); - } - - // Print labels, that is, GPU noise - if (!result.report_label.empty()) - { - std::string s = FormatString("%.2f", 100. * std::stod(result.report_label)); - const char* unit = "%"; - PrintColoredString(sout, get_log_color(COLOR_DEFAULT), " %13s%s", s.c_str(), unit); - } - - PrintColoredString(sout, get_log_color(COLOR_DEFAULT), "\n"); - } - - // Called once for each group of benchmark runs, gives information about - // cpu-time, gpu-time, elements processed, global memory bw and the % - // the latter represents from the peak global bandwidth during the - // benchmark run. If the group of runs contained more than two entries - // then 'report' contains additional elements representing the mean, - // standard deviation and coefficient of variation of those runs. - // Additionally if this group of runs was the last in a family of - // benchmarks, 'reports' contains additional entries representing the - // asymptotic complexity and RMS of that benchmark family. - void ReportRuns(const std::vector& reports) - { - for (const auto& run : reports) - { - // Print the header if none was printed yet - bool print_header = !printed_header_; - if (print_header) - { - printed_header_ = true; - PrintHeader(run); - } - PrintRunData(run); - } - } -}; - -class CustomJSONReporter : public benchmark::JSONReporter -{ -private: - bool first_report_ = true; - - std::string StrEscape(const std::string& s) - { - std::string tmp; - tmp.reserve(s.size()); - for (char c : s) - { - switch (c) - { - case '\b': - tmp += "\\b"; - break; - case '\f': - tmp += "\\f"; - break; - case '\n': - tmp += "\\n"; - break; - case '\r': - tmp += "\\r"; - break; - case '\t': - tmp += "\\t"; - break; - case '\\': - tmp += "\\\\"; - break; - case '"': - tmp += "\\\""; - break; - default: - tmp += c; - break; - } - } - return tmp; - } - - std::string FormatKV(std::string const& key, std::string const& value) - { - return FormatString("\"%s\": \"%s\"", StrEscape(key).c_str(), StrEscape(value).c_str()); - } - - std::string FormatKV(std::string const& key, const char* value) - { - return FormatString("\"%s\": \"%s\"", StrEscape(key).c_str(), StrEscape(value).c_str()); - } - - std::string FormatKV(std::string const& key, bool value) - { - return FormatString("\"%s\": %s", StrEscape(key).c_str(), value ? "true" : "false"); - } - - std::string FormatKV(std::string const& key, int64_t value) - { - std::stringstream ss; - ss << '"' << StrEscape(key) << "\": " << value; - return ss.str(); - } - - std::string FormatKV(std::string const& key, double value) - { - std::stringstream ss; - ss << '"' << StrEscape(key) << "\": "; - - if (std::isnan(value)) - { - ss << (value < 0 ? "-" : "") << "NaN"; - } - else if (std::isinf(value)) - { - ss << (value < 0 ? "-" : "") << "Infinity"; - } - else - { - const auto max_digits10 = std::numeric_limits::max_digits10; - const auto max_fractional_digits10 = max_digits10 - 1; - ss << std::scientific << std::setprecision(max_fractional_digits10) << value; - } - return ss.str(); - } - -public: - void ReportRuns(std::vector const& reports) - { - if (reports.empty()) - { - return; - } - std::string indent(4, ' '); - std::ostream& out = GetOutputStream(); - if (!first_report_) - { - out << ",\n"; - } - first_report_ = false; - - for (auto it = reports.begin(); it != reports.end(); ++it) - { - out << indent << "{\n"; - PrintRunData(*it); - out << indent << '}'; - auto it_cp = it; - if (++it_cp != reports.end()) - { - out << ",\n"; - } - } - } - - void PrintRunData(Run const& run) - { - std::string indent(6, ' '); - std::ostream& out = GetOutputStream(); - - auto output_format = [this, &out, &indent](const std::string& label, auto val, bool start_endl = true) { - if (start_endl) - { - out << ",\n"; - } - out << indent << FormatKV(label, val); - }; - - output_format("name", run.benchmark_name(), false); - output_format("family_index", run.family_index); - output_format("per_family_instance_index", run.per_family_instance_index); - output_format("run_name", run.run_name.str()); - output_format("run_type", [&run]() -> const char* { - switch (run.run_type) - { - case BenchmarkReporter::Run::RT_Iteration: - return "iteration"; - case BenchmarkReporter::Run::RT_Aggregate: - return "aggregate"; - } - BENCHMARK_UNREACHABLE(); - }()); - output_format("repetitions", run.repetitions); - if (run.run_type != BenchmarkReporter::Run::RT_Aggregate) - { - output_format("repetition_index", run.repetition_index); - } - output_format("threads", run.threads); - if (run.run_type == BenchmarkReporter::Run::RT_Aggregate) - { - output_format("aggregate_name", run.aggregate_name); - output_format("aggregate_unit", [&run]() -> const char* { - switch (run.aggregate_unit) - { - case benchmark::StatisticUnit::kTime: - return "time"; - case benchmark::StatisticUnit::kPercentage: - return "percentage"; - } - BENCHMARK_UNREACHABLE(); - }()); - } - if (benchmark::internal::SkippedWithError == run.skipped) - { - output_format("error_occurred", true); - output_format("error_message", run.skip_message); - } - else if (benchmark::internal::SkippedWithMessage == run.skipped) - { - output_format("skipped", true); - output_format("skip_message", run.skip_message); - } - if (!run.report_big_o && !run.report_rms) - { - output_format("iterations", run.iterations); - if (run.run_type != Run::RT_Aggregate || run.aggregate_unit == benchmark::StatisticUnit::kTime) - { - output_format("gpu_time", run.GetAdjustedRealTime()); - output_format("cpu_time", run.GetAdjustedCPUTime()); - } - else - { - assert(run.aggregate_unit == benchmark::StatisticUnit::kPercentage); - output_format("gpu_time", run.real_accumulated_time); - output_format("cpu_time", run.cpu_accumulated_time); - } - output_format("time_unit", GetTimeUnitString(run.time_unit)); - } - else if (run.report_big_o) - { - output_format("cpu_coefficient", run.GetAdjustedCPUTime()); - output_format("gpu_coefficient", run.GetAdjustedRealTime()); - output_format("big_o", get_complexity(run.complexity)); - output_format("time_unit", GetTimeUnitString(run.time_unit)); - } - else if (run.report_rms) - { - output_format("rms", run.GetAdjustedCPUTime()); - } - - for (auto& c : run.counters) - { - if (c.first == "items_per_second") - { - // Report same name as console reporter - output_format("elements_per_second", c.second); - } - else if (c.first == "bytes_per_second") - { - // Report same name as console reporter - output_format("global_mem_bw", c.second); - const double util_bw = calculate_bw_utils(run); - if (util_bw >= 0) - { - output_format("util_bw", util_bw); - } - } - else - { - output_format(c.first, c.second); - } - } - - const benchmark::MemoryManager::Result memory_result = run.memory_result; - output_format("allocs_per_iter", run.allocs_per_iter); - output_format("max_bytes_used", memory_result.max_bytes_used); - - if (memory_result.total_allocated_bytes != benchmark::MemoryManager::TombstoneValue) - { - output_format("total_allocated_bytes", memory_result.total_allocated_bytes); - } - - if (memory_result.net_heap_growth != benchmark::MemoryManager::TombstoneValue) - { - output_format("net_heap_growth", memory_result.net_heap_growth); - } - - if (!run.report_label.empty()) - { - output_format("label", run.report_label); - } - out << '\n'; - } -}; - -BENCHMARK_DISABLE_DEPRECATED_WARNING - -class CustomCSVReporter : public benchmark::CSVReporter -{ -private: - bool printed_header_ = false; - - std::vector elements = { - "name", - "iterations", - "gpu_time", - "cpu_time", - "time_unit", - "global_mem_bw", - "util_bw", - "elements_per_second", - "gpu_noise", - "label", - "error_occurred", - "error_message"}; - - std::string CsvEscape(const std::string& s) - { - std::string tmp; - tmp.reserve(s.size() + 2); - for (char c : s) - { - switch (c) - { - case '"': - tmp += "\"\""; - break; - default: - tmp += c; - break; - } - } - return '"' + tmp + '"'; - } - -public: - void PrintHeader(const Run& /*run*/) - { - std::string str = ""; - bool first = true; - for (auto element : elements) - { - if (first) - { - first = false; - } - else - { - str += ","; - } - str += element; - } - GetOutputStream() << str << "\n"; - } - - void PrintRunData(const Run& result) - { - // Report benchmark name - auto& sout = GetOutputStream(); - - if (result.skipped) - { - sout << std::string(elements.size() - 3, ','); - sout << std::boolalpha << (benchmark::internal::SkippedWithError == result.skipped) << ","; - sout << CsvEscape(result.skip_message) << "\n"; - return; - } - - sout << CsvEscape(result.benchmark_name()) << ","; - - if (!result.report_big_o && !result.report_rms) - { - sout << result.iterations; - } - sout << ","; - - sout << result.GetAdjustedRealTime() << ","; - sout << result.GetAdjustedCPUTime() << ","; - - if (result.report_big_o) - { - sout << get_complexity(result.complexity); - } - else if (!result.report_rms) - { - sout << benchmark::GetTimeUnitString(result.time_unit); - } - sout << ","; - - if (result.counters.find("bytes_per_second") != result.counters.end()) - { - sout << result.counters.at("bytes_per_second"); - } - sout << ","; - - if (result.counters.find("bytes_per_second") != result.counters.end()) - { - const double bw_util = calculate_bw_utils(result); - if (bw_util >= 0) - { - sout << bw_util; - } - } - sout << ","; - - if (result.counters.find("items_per_second") != result.counters.end()) - { - sout << result.counters.at("items_per_second"); - } - sout << ","; - - if (result.counters.find("gpu_noise") != result.counters.end()) - { - sout << result.counters.at("gpu_noise"); - } - sout << ","; - - if (!result.report_label.empty()) - { - sout << CsvEscape(result.report_label); - } - - sout << ",,\n"; - } - - void ReportRuns(const std::vector& reports) - { - for (const auto& run : reports) - { - // Print the header if none was printed yet - bool print_header = !printed_header_; - if (print_header) - { - printed_header_ = true; - PrintHeader(run); - } - PrintRunData(run); - } - } -}; - -benchmark::BenchmarkReporter* ChooseCustomReporter() -{ - // benchmark::BenchmarkReporter is polymorphic as it has a virtual - // function which allows us to use dynamic_cast to detect the derived type. - using PtrType = benchmark::BenchmarkReporter*; - PtrType default_display_reporter = benchmark::CreateDefaultDisplayReporter(); - - if (IsType(default_display_reporter)) - { - return PtrType(new CustomCSVReporter); - } - else if (IsType(default_display_reporter)) - { - return PtrType(new CustomJSONReporter); - } - else if (IsType(default_display_reporter)) - { - return PtrType(new CustomConsoleReporter); - } - - return nullptr; -} - -BENCHMARK_RESTORE_DEPRECATED_WARNING - -} // namespace bench_utils -#endif // ROCTHRUST_BENCHMARKS_BENCH_UTILS_CUSTOM_REPORTER_HPP_ diff --git a/projects/rocthrust/benchmark/bench_utils/generation_utils.hpp b/projects/rocthrust/benchmark/generation_utils.hpp similarity index 71% rename from projects/rocthrust/benchmark/bench_utils/generation_utils.hpp rename to projects/rocthrust/benchmark/generation_utils.hpp index 607f36016343..0d163af8f4fe 100644 --- a/projects/rocthrust/benchmark/bench_utils/generation_utils.hpp +++ b/projects/rocthrust/benchmark/generation_utils.hpp @@ -1,6 +1,6 @@ // MIT License // -// Copyright (c) 2024-2025 Advanced Micro Devices, Inc. All rights reserved. +// Copyright (c) 2024-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,11 +20,7 @@ // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE // SOFTWARE. -#ifndef ROCTHRUST_BENCHMARKS_BENCH_UTILS_GENERATION_UTILS_HPP_ -#define ROCTHRUST_BENCHMARKS_BENCH_UTILS_GENERATION_UTILS_HPP_ - -// Utils -#include "common/types.hpp" +#pragma once // Thrust #include @@ -46,15 +42,9 @@ #include #include -// rocPRIM -#include - // rocRAND #include -// Google Benchmark -#include - // STL #include #include @@ -75,45 +65,6 @@ namespace bench_utils { -/// \brief Provides a sequence of seeds. -class managed_seed -{ -public: - /// \param[in] seed_string Either "random" to get random seeds, - /// or an unsigned integer to get (a sequence) of deterministic seeds. - managed_seed(const std::string& seed_string) - { - is_random = seed_string == "random"; - if (!is_random) - { - const unsigned int seed = std::stoul(seed_string); - std::seed_seq seq{seed}; - seq.generate(seeds.begin(), seeds.end()); - } - } - - managed_seed() - : managed_seed("random"){}; - - unsigned int get_0() const - { - return is_random ? std::random_device{}() : seeds[0]; - } - - unsigned int get_1() const - { - return is_random ? std::random_device{}() : seeds[1]; - } - - unsigned int get_2() const - { - return is_random ? std::random_device{}() : seeds[2]; - } - -private: - std::array seeds; - bool is_random; -}; float get_entropy_percentage(int entropy_reduction) { @@ -135,7 +86,7 @@ float get_entropy_percentage(int entropy_reduction) } template -T value_from_entropy(float64_t percentage) +T value_from_entropy(double percentage) { if (percentage == 100.0) { @@ -189,11 +140,22 @@ struct random_to_item_t::va } }; +template +struct geq_t +{ + T val; + + __host__ __device__ bool operator()(T x) + { + return x >= val; + } +}; + // Integral types template struct random_to_item_t::value>::type> { -#if THRUST_BENCHMARKS_HAVE_INT128_SUPPORT +#ifndef _MSC_VER using CastT = typename std::conditional::value || std::is_same::value, typename std::conditional::value, long, unsigned long>::type, @@ -237,65 +199,20 @@ struct and_t } }; -template -struct geq_t -{ - T val; - - __host__ __device__ bool operator()(T x) - { - return x >= val; - } -}; - -template -class value_wrapper_t -{ - T m_val{}; - -public: - explicit value_wrapper_t(T val) - : m_val(val) - {} - - T get() const - { - return m_val; - } - - value_wrapper_t& operator++() - { - m_val++; - return *this; - } -}; - -class seed_t : public value_wrapper_t -{ -public: - using value_wrapper_t::value_wrapper_t; - using value_wrapper_t::operator++; - - seed_t() - : value_wrapper_t(42) - {} -}; - struct device_generator_base_t { const std::size_t elements{0}; - const std::string seed_type{"random"}; - seed_t seed{}; + const uint32_t seed_source{}; + uint32_t seed{}; const int entropy_reduction{0 /*bit_entropy::_1_000*/}; - device_generator_base_t(std::size_t m_elements, const std::string& m_seed_type, int m_entropy_reduction) + device_generator_base_t(std::size_t m_elements, const uint32_t& m_seed_source, int m_entropy_reduction) : elements(m_elements) - , seed_type(m_seed_type) + , seed_source(m_seed_source) , entropy_reduction(m_entropy_reduction) { ROCRAND_CHECK(rocrand_create_generator(&gen, ROCRAND_RNG_PSEUDO_DEFAULT)); - const managed_seed managed_seed{seed_type}; - seed = seed_t{managed_seed.get_0()}; + seed = seed_source; } ~device_generator_base_t() @@ -320,7 +237,7 @@ struct device_generator_base_t else if (entropy_reduction >= 5) /*bit_entropy::_0_000*/ { std::mt19937 rng; - rng.seed(static_cast(seed.get())); + rng.seed(static_cast(seed)); std::uniform_real_distribution dist(0.0f, 1.0f); T random_value = detail::random_to_item_t(min, max)(dist(rng)); thrust::fill(policy, data.data(), data.data() + data.size(), random_value); @@ -348,12 +265,12 @@ struct device_generator_base_t } } - const double* new_uniform_distribution(seed_t seed, std::size_t num_items) + const double* new_uniform_distribution(uint32_t seed, std::size_t num_items) { distribution.resize(num_items); double* d_distribution = thrust::raw_pointer_cast(distribution.data()); - ROCRAND_CHECK(rocrand_set_seed(gen, seed.get())); + ROCRAND_CHECK(rocrand_set_seed(gen, seed)); ROCRAND_CHECK(rocrand_generate_uniform_double(gen, d_distribution, num_items)); hipError_t error = hipDeviceSynchronize(); @@ -393,8 +310,8 @@ struct device_vector_generator_t : device_generator_base_t const T max{std::numeric_limits::max()}; device_vector_generator_t( - std::size_t m_elements, const std::string& m_seed_type, int m_entropy_reduction, T m_min, T m_max) - : device_generator_base_t(m_elements, m_seed_type, m_entropy_reduction) + std::size_t m_elements, const uint32_t& m_seed_source, int m_entropy_reduction, T m_min, T m_max) + : device_generator_base_t(m_elements, m_seed_source, m_entropy_reduction) , min(m_min) , max(m_max) {} @@ -408,8 +325,8 @@ struct device_vector_generator_t : device_generator_base_t template <> struct device_vector_generator_t : device_generator_base_t { - device_vector_generator_t(std::size_t m_elements, const std::string& m_seed_type, int m_entropy_reduction) - : device_generator_base_t(m_elements, m_seed_type, m_entropy_reduction) + device_vector_generator_t(std::size_t m_elements, const uint32_t& m_seed_source, int m_entropy_reduction) + : device_generator_base_t(m_elements, m_seed_source, m_entropy_reduction) {} template @@ -421,14 +338,14 @@ struct device_vector_generator_t : device_generator_base_t template std::size_t gen_uniform_offsets( - const std::string seed_type, + const uint32_t& seed_source, thrust::device_vector& segment_offsets, const std::size_t min_segment_size, const std::size_t max_segment_size) { const T elements = segment_offsets.size() - 2; - segment_offsets = device_generator_base_t(segment_offsets.size(), seed_type, 0 /*bit_entropy::_1_000*/) + segment_offsets = device_generator_base_t(segment_offsets.size(), seed_source, 0 /*bit_entropy::_1_000*/) .generate(static_cast(min_segment_size), static_cast(max_segment_size)); // Find the range of contiguous offsets starting from index 0 which sum is greater or @@ -501,90 +418,20 @@ void gen_key_segments(thrust::device_vector& keys, thrust::device_vector -// struct repeat_index_t -// { -// __host__ __device__ __forceinline__ thrust::constant_iterator operator()(std::size_t i) -// { -// return thrust::constant_iterator(static_cast(i)); -// } -// }; - -// template -// struct offset_to_iterator_t -// { -// T* base_it; - -// __host__ __device__ __forceinline__ T* operator()(std::size_t offset) const -// { -// return base_it + offset; -// } -// }; - -// struct offset_to_size_t -// { -// std::size_t* offsets = nullptr; - -// __host__ __device__ __forceinline__ std::size_t operator()(std::size_t i) -// { -// return offsets[i + 1] - offsets[i]; -// } -// }; -// -// template -// void gen_key_segments(thrust::device_vector& keys, -// thrust::device_vector& segment_offsets) -// { - -// const std::size_t total_segments = segment_offsets.size() - 1; - -// thrust::counting_iterator iota(0); -// repeat_index_t src_transform_op {}; -// offset_to_iterator_t dst_transform_op {thrust::raw_pointer_cast(keys.data())}; -// offset_to_size_t size_transform_op {thrust::raw_pointer_cast(segment_offsets.data())}; - -// auto d_range_srcs = thrust::make_transform_iterator(iota, src_transform_op); -// auto d_range_dsts -// = thrust::make_transform_iterator(segment_offsets.begin(), dst_transform_op); -// auto d_range_sizes = thrust::make_transform_iterator(iota, size_transform_op); - -// std::uint8_t* d_temp_storage = nullptr; -// std::size_t temp_storage_bytes = 0; - -// rocprim::batch_copy(d_temp_storage, -// temp_storage_bytes, -// d_range_srcs, -// d_range_dsts, -// d_range_sizes, -// total_segments); - -// thrust::device_vector temp_storage(temp_storage_bytes); -// d_temp_storage = thrust::raw_pointer_cast(temp_storage.data()); - -// rocprim::batch_copy(d_temp_storage, -// temp_storage_bytes, -// d_range_srcs, -// d_range_dsts, -// d_range_sizes, -// total_segments); -// hipDeviceSynchronize(); -// } - struct device_uniform_key_segments_generator_t { const std::size_t elements{0}; - const std::string seed_type{"random"}; + const uint32_t seed_source{}; const std::size_t min_segment_size{0}; const std::size_t max_segment_size{0}; device_uniform_key_segments_generator_t( std::size_t m_elements, - const std::string m_seed_type, + const uint32_t& m_seed_source, const std::size_t m_min_segment_size, const std::size_t m_max_segment_size) : elements(m_elements) - , seed_type(m_seed_type) + , seed_source(m_seed_source) , min_segment_size(m_min_segment_size) , max_segment_size(m_max_segment_size) {} @@ -596,7 +443,7 @@ struct device_uniform_key_segments_generator_t thrust::device_vector segment_offsets(keys.size() + 2); const std::size_t offsets_size = - gen_uniform_offsets(seed_type, segment_offsets, min_segment_size, max_segment_size); + gen_uniform_offsets(seed_source, segment_offsets, min_segment_size, max_segment_size); segment_offsets.resize(offsets_size); gen_key_segments(keys, segment_offsets); @@ -609,11 +456,11 @@ struct gen_uniform_key_segments_t { device_uniform_key_segments_generator_t operator()( const std::size_t elements, - const std::string seed_type, + const uint32_t& seed_source, const std::size_t min_segment_size, const std::size_t max_segment_size) const { - return {elements, seed_type, min_segment_size, max_segment_size}; + return {elements, seed_source, min_segment_size, max_segment_size}; } }; @@ -627,18 +474,18 @@ struct gen_t template device_vector_generator_t operator()( std::size_t elements, - const std::string seed_type, + const uint32_t& seed_source, const int entropy = 0 /*100*/, T min = std::numeric_limits::min, T max = std::numeric_limits::max()) const { - return {elements, seed_type, entropy, min, max}; + return {elements, seed_source, entropy, min, max}; } device_vector_generator_t - operator()(std::size_t elements, const std::string seed_type, const int entropy = 0 /*100*/) const + operator()(std::size_t elements, const uint32_t& seed_source, const int entropy = 0 /*100*/) const { - return {elements, seed_type, entropy}; + return {elements, seed_source, entropy}; } gen_uniform_t uniform{}; @@ -648,4 +495,3 @@ struct gen_t detail::gen_t generate; } // namespace bench_utils -#endif // ROCTHRUST_BENCHMARKS_BENCH_UTILS_GENERATION_UTILS_HPP_ From 529e814482ae4e3af5e1a780c1e5ea52a6b2a207 Mon Sep 17 00:00:00 2001 From: NguyenNhuDi Date: Fri, 17 Jul 2026 11:40:14 -0400 Subject: [PATCH 03/12] feat(rocthrust): migrate adjacent_difference to primbench --- .../bench/adjacent_difference/basic.cu | 154 +++++----------- .../bench/adjacent_difference/custom.cu | 169 +++++------------- .../bench/adjacent_difference/in_place.cu | 150 +++++----------- 3 files changed, 131 insertions(+), 342 deletions(-) diff --git a/projects/rocthrust/benchmark/bench/adjacent_difference/basic.cu b/projects/rocthrust/benchmark/bench/adjacent_difference/basic.cu index 33f8f023449b..3a72b927f675 100644 --- a/projects/rocthrust/benchmark/bench/adjacent_difference/basic.cu +++ b/projects/rocthrust/benchmark/bench/adjacent_difference/basic.cu @@ -1,6 +1,6 @@ /****************************************************************************** * Copyright (c) 2011-2023, NVIDIA CORPORATION. All rights reserved. - * Modifications Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved. + * Modifications Copyright (c) 2024-2026, Advanced Micro Devices, Inc. All rights reserved. * * Redistribution and use in source and binary forms, with or without * modification, are permitted provided that the following conditions are met: @@ -26,135 +26,69 @@ * ******************************************************************************/ -// Benchmark utils -#include "../../bench_utils/bench_utils.hpp" - -// rocThrust #include #include #include -// Google Benchmark -#include - -// STL -#include -#include -#include +#include "bench_utils.hpp" -struct basic +template +struct adjacent_difference_benchmark : public primbench::benchmark_interface { - template - float64_t run(thrust::device_vector& input, thrust::device_vector& output, Policy policy) - { - thrust::adjacent_difference(policy, input.cbegin(), input.cend(), output.begin()); + adjacent_difference_benchmark(size_t items) + : m_items(items) + {} - bench_utils::gpu_timer d_timer; - - d_timer.start(0); - thrust::adjacent_difference(policy, input.cbegin(), input.cend(), output.begin()); - d_timer.stop(0); - - return d_timer.get_duration(); + primbench::json meta() const override + { + return primbench::json{} + .add("algo", "adjacent_difference") + .add("subalgo", "basic") + .add("input_type", primbench::name()) + .add("elements", m_items); } -}; - -template -void run_benchmark(benchmark::State& state, const std::size_t elements, const std::string seed_type) -{ - // Benchmark object - Benchmark benchmark{}; - - // GPU times - std::vector gpu_times; - - // Generate input - thrust::device_vector input = bench_utils::generate(elements, seed_type); - - // Output - thrust::device_vector output(elements); - bench_utils::caching_allocator_t alloc{}; - thrust::detail::device_t policy{}; - - for (auto _ : state) + void run(primbench::state& state) override { - float64_t duration = benchmark.template run(input, output, policy(alloc)); - state.SetIterationTime(duration); - gpu_times.push_back(duration); - } + bench_utils::caching_allocator_t alloc{}; + thrust::detail::device_t policy{}; - // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported - // it will actually be the global memory bandwidth gotten. - state.SetBytesProcessed(state.iterations() * 2 * elements * sizeof(T)); - state.SetItemsProcessed(state.iterations() * elements); + thrust::device_vector in = bench_utils::generate(m_items, state.seed); - const double gpu_cv = bench_utils::StatisticsCV(gpu_times); - state.counters["gpu_noise"] = gpu_cv; -} + thrust::device_vector out(m_items); -#define CREATE_BENCHMARK(T, Elements) \ - benchmark::RegisterBenchmark( \ - bench_utils::bench_naming::format_name("{algo:adjacent_difference,subalgo:" + name + ",input_type:" #T \ - + ",elements:" + bench_utils::format_pow2(Elements)) \ - .c_str(), \ - run_benchmark, \ - Elements, \ - seed_type) + state.set_items(m_items); + state.add_reads(m_items); + state.add_writes(m_items); -#define BENCHMARK_TYPE(type) \ - for (size_t size : bench_utils::sizes) \ - { \ - if (sizeof(type) * size <= bench_utils::system.devProp.totalGlobalMem) \ - bs.push_back(CREATE_BENCHMARK(type, size)); \ + state.run([&] { + thrust::adjacent_difference(policy(alloc).on(state.stream), in.cbegin(), in.cend(), out.begin()); + }); } -template -void add_benchmarks( - const std::string& name, std::vector& benchmarks, const std::string seed_type) -{ - std::vector bs; - BENCHMARK_TYPE(int8_t) - BENCHMARK_TYPE(int16_t) - BENCHMARK_TYPE(int32_t) - BENCHMARK_TYPE(int64_t) - BENCHMARK_TYPE(float32_t) - BENCHMARK_TYPE(float64_t) - benchmarks.insert(benchmarks.end(), bs.begin(), bs.end()); -} +private: + size_t m_items; +}; + +#define QUEUE(T) \ + for (size_t size : bench_utils::sizes(2 * sizeof(T))) \ + executor.queue>(size); int main(int argc, char* argv[]) { - cli::Parser parser(argc, argv); - parser.set_optional("name_format", "name_format", "human", "either: json,human,txt"); - parser.set_optional("seed", "seed", "random", bench_utils::get_seed_message()); - parser.run_and_exit_if_error(); - - // Parse argv - benchmark::Initialize(&argc, argv); - bench_utils::bench_naming::set_format(parser.get("name_format")); /* either: json,human,txt */ - const std::string seed_type = parser.get("seed"); - - // Benchmark info - bench_utils::add_common_benchmark_info(); - benchmark::AddCustomContext("seed", seed_type); + primbench::settings settings; + settings.size = 1; // bench_utils::sizes() calculates it later. + settings.min_gpu_ms_per_batch = 100; + settings.batch_window_size = 3; + primbench::executor executor(argc, argv, settings, primbench::flags::sync); - // Add benchmark - std::vector benchmarks; - add_benchmarks("basic", benchmarks, seed_type); - - // Use manual timing - for (auto& b : benchmarks) - { - b->UseManualTime(); - b->Unit(benchmark::kMicrosecond); - b->MinTime(0.4); // in seconds - } + QUEUE(int8_t) + QUEUE(int16_t) + QUEUE(int32_t) + QUEUE(int64_t) - // Run benchmarks - benchmark::RunSpecifiedBenchmarks(bench_utils::ChooseCustomReporter()); + QUEUE(float) + QUEUE(double) - // Finish - benchmark::Shutdown(); - return 0; + executor.run(); } diff --git a/projects/rocthrust/benchmark/bench/adjacent_difference/custom.cu b/projects/rocthrust/benchmark/bench/adjacent_difference/custom.cu index 2be6dc85c80e..50793dde5c36 100644 --- a/projects/rocthrust/benchmark/bench/adjacent_difference/custom.cu +++ b/projects/rocthrust/benchmark/bench/adjacent_difference/custom.cu @@ -1,6 +1,6 @@ /****************************************************************************** * Copyright (c) 2011-2023, NVIDIA CORPORATION. All rights reserved. - * Modifications Copyright (c) 2024-2025, Advanced Micro Devices, Inc. All rights reserved. + * Modifications Copyright (c) 2024-2026, Advanced Micro Devices, Inc. All rights reserved. * * Redistribution and use in source and binary forms, with or without * modification, are permitted provided that the following conditions are met: @@ -26,154 +26,75 @@ * ******************************************************************************/ -// Benchmark utils -#include "../../bench_utils/bench_utils.hpp" - -// rocThrust #include #include #include -// Google Benchmark -#include +#include "bench_utils.hpp" -// STL -#include -#include -#include +#define VAL 42 template -struct custom_op +struct adjacent_difference_benchmark : public primbench::benchmark_interface { - T val; - - custom_op() = delete; - - explicit custom_op(T val) - : val(val) + adjacent_difference_benchmark(size_t items) + : m_items(items) {} - __device__ T operator()(const T& lhs, const T& rhs) + primbench::json meta() const override { - return lhs * rhs + val; // Hope to gen mad + return primbench::json{} + .add("algo", "adjacent_difference") + .add("subalgo", "custom") + .add("input_type", primbench::name()) + .add("elements", m_items); } -}; -template -struct custom -{ - template - float64_t run(thrust::device_vector& input, thrust::device_vector& output, Policy policy) + void run(primbench::state& state) override { - thrust::adjacent_difference(policy, input.cbegin(), input.cend(), output.begin(), custom_op{Val}); - - bench_utils::gpu_timer d_timer; - - d_timer.start(0); - thrust::adjacent_difference(policy, input.cbegin(), input.cend(), output.begin(), custom_op{Val}); - d_timer.stop(0); - - return d_timer.get_duration(); - } -}; - -template -void run_benchmark(benchmark::State& state, const std::size_t elements, const std::string seed_type) -{ - // Benchmark object - Benchmark benchmark{}; + auto custom_op = [] __device__(const T& lhs, const T& rhs) { + return lhs * rhs + VAL; + }; - // GPU times - std::vector gpu_times; + bench_utils::caching_allocator_t alloc{}; + thrust::detail::device_t policy{}; - // Generate input - thrust::device_vector input = bench_utils::generate(elements, seed_type); + thrust::device_vector in = bench_utils::generate(m_items, state.seed); - // Output - thrust::device_vector output(elements); + thrust::device_vector out(m_items); - bench_utils::caching_allocator_t alloc{}; - thrust::detail::device_t policy{}; + state.set_items(m_items); + state.add_reads(m_items); + state.add_writes(m_items); - for (auto _ : state) - { - float64_t duration = benchmark.template run(input, output, policy(alloc)); - state.SetIterationTime(duration); - gpu_times.push_back(duration); + state.run([&] { + thrust::adjacent_difference(policy(alloc).on(state.stream), in.cbegin(), in.cend(), out.begin(), custom_op); + }); } - // BytesProcessed include read and written bytes, so when the BytesProcessed/s are reported - // it will actually be the global memory bandwidth gotten. - state.SetBytesProcessed(state.iterations() * 2 * elements * sizeof(T)); - state.SetItemsProcessed(state.iterations() * elements); - - const double gpu_cv = bench_utils::StatisticsCV(gpu_times); - state.counters["gpu_noise"] = gpu_cv; -} +private: + size_t m_items; +}; -#define CREATE_BENCHMARK(T, Elements) \ - benchmark::RegisterBenchmark( \ - bench_utils::bench_naming::format_name("{algo:adjacent_difference,subalgo:" + name + ",input_type:" #T \ - + ",elements:" + bench_utils::format_pow2(Elements)) \ - .c_str(), \ - run_benchmark, T>, \ - Elements, \ - seed_type) - -#define BENCHMARK_TYPE(type) \ - for (size_t size : bench_utils::sizes) \ - { \ - if (sizeof(type) * size <= bench_utils::system.devProp.totalGlobalMem) \ - bs.push_back(CREATE_BENCHMARK(type, size)); \ - } - -template