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/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/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/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