diff --git a/projects/rocprim/.gitlab/run_benchmarks.py b/projects/rocprim/.gitlab/run_benchmarks.py index f07b7e11f406..d456de43e013 100755 --- a/projects/rocprim/.gitlab/run_benchmarks.py +++ b/projects/rocprim/.gitlab/run_benchmarks.py @@ -50,8 +50,6 @@ def run_benchmarks( max_gpu_temp, max_warming_secs, max_cooling_secs, - output_hip_device_properties_context, - output_amdsmi_context, output_batches, spaces_per_indent, stream_blocking_timeout_secs, @@ -148,10 +146,6 @@ def should_skip(json_out_path): args += ["--max-warming-secs", max_warming_secs] if max_cooling_secs: args += ["--max-cooling-secs", max_cooling_secs] - if output_hip_device_properties_context: - args += ["--output-hip-device-properties-context"] - if output_amdsmi_context: - args += ["--output-amdsmi-context"] if output_batches: args += ["--output-batches"] if spaces_per_indent: @@ -294,20 +288,6 @@ def main(): default="", required=False, ) - parser.add_argument( - "--output-hip-device-properties-context", - help="Output a `hip_device_properties` object in the context, containing GPU details", - default=False, - action="store_true", - required=False, - ) - parser.add_argument( - "--output-amdsmi-context", - help="Output an `amdsmi` object in the context, containing GPU details", - default=False, - action="store_true", - required=False, - ) parser.add_argument( "--output-batches", help="Output a `batches` array for each specialization, containing per-batch details", @@ -351,8 +331,6 @@ def main(): args.max_gpu_temp, args.max_warming_secs, args.max_cooling_secs, - args.output_hip_device_properties_context, - args.output_amdsmi_context, args.output_batches, args.spaces_per_indent, args.stream_blocking_timeout_secs, diff --git a/shared/primbench/README.md b/shared/primbench/README.md index 6a9006be4739..ecf62d4c5ab6 100644 --- a/shared/primbench/README.md +++ b/shared/primbench/README.md @@ -1,6 +1,6 @@ # primbench -primbench is a single-header HIP benchmarking library. +primbench is a single-header HIP and CUDA benchmarking library. ![primbench demo](./assets/primbench.gif) @@ -16,8 +16,8 @@ primbench is a single-header HIP benchmarking library. ## Dependencies primbench has the following dependencies: -- HIP -- AMD SMI (for querying live GPU statistics) +- HIP or CUDA +- [AMD SMI](https://rocm.docs.amd.com/projects/amdsmi/en/latest/) or [NVML](https://developer.nvidia.com/management-library-nvml) (for querying live GPU statistics) - C++17 or later ## Example @@ -81,15 +81,15 @@ struct copy_benchmark : public primbench::benchmark_interface primbench::log("Allocating device memory"); T* d_input; T* d_output; - PRIMBENCH_HIP_CHECK(hipMalloc(&d_input, items * sizeof(T))); - PRIMBENCH_HIP_CHECK(hipMalloc(&d_output, items * sizeof(T))); + PRIMBENCH_CHECK(hipMalloc(&d_input, items * sizeof(T))); + PRIMBENCH_CHECK(hipMalloc(&d_output, items * sizeof(T))); primbench::log("Copying to device"); - PRIMBENCH_HIP_CHECK(hipMemcpyAsync(d_input, - h_input.data(), - items * sizeof(T), - hipMemcpyHostToDevice, - stream)); + PRIMBENCH_CHECK(hipMemcpyAsync(d_input, + h_input.data(), + items * sizeof(T), + hipMemcpyHostToDevice, + stream)); dim3 grid(items / items_per_block); dim3 block(BlockSize); @@ -107,8 +107,8 @@ struct copy_benchmark : public primbench::benchmark_interface <<>>(d_input, d_output); }); - PRIMBENCH_HIP_CHECK(hipFree(d_input)); - PRIMBENCH_HIP_CHECK(hipFree(d_output)); + PRIMBENCH_CHECK(hipFree(d_input)); + PRIMBENCH_CHECK(hipFree(d_output)); } }; @@ -123,29 +123,49 @@ int main(int argc, char* argv[]) } ``` -The benchmark is compiled and run like so: +The HIP benchmark is compiled and run like so: ```bash -hipcc -o copy_benchmark examples/copy_benchmark.cpp -I. -lamd_smi && ./copy_benchmark +hipcc -o copy_benchmark examples/hip/copy_benchmark.cpp -I. -lamd_smi && ./copy_benchmark +``` + +And its equivalent CUDA benchmark is compiled and run like so: + +```bash +nvcc -o copy_benchmark examples/cuda/copy_benchmark.cu -I. -lnvidia-ml && ./copy_benchmark ``` It outputs this `results.json`: ```json { "context": { - "results_version": "2.0.0", + "results_version": "3.0.0", "general": { "algorithm": "copy", "specialization_count": 2, - "gpu_name": "AMD Instinct MI210", - "gpu_arch": "gfx90a", - "gpu_pci_bus_id": "0000:83:00.0", "library_build_type": "debug", - "temp_type": "edge", + "gpu": { + "name": "AMD Radeon RX 9070 XT", + "arch": "gfx1201", + "pci_bus_id": "0000:83:00.0" + }, + "backend": { + "name": "hip", + "hip_version": "6.4.43482-0f2d60242", + "runtime_version": "6.4.43482", + "driver_version": "6.4.43482", + "compiler": { + "name": "clang", + "version": "19.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.4.0 25133 c7fe45cf4b819c5991fe208aaa96edf142730f1d)" + }, + "monitoring": { + "name": "amdsmi", + "version": "25.3.0" + } + }, + "temperature_type": "edge", "host_name": "host", - "date": "2025-11-24T15:01:50+00:00", - "hip_version": "6.4.43482-0f2d60242", - "clang_version": "19.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.4.0 25133 c7fe45cf4b819c5991fe208aaa96edf142730f1d)" + "date": "2025-11-24T15:01:50+00:00" }, "settings": { "size": 134217728, @@ -164,8 +184,6 @@ It outputs this `results.json`: "max_gpu_temp": 60, "max_warming_secs": 60, "max_cooling_secs": 60, - "output_hip_device_properties_context": false, - "output_amdsmi_context": false, "output_batches": false, "spaces_per_indent": 4, "stream_blocking_timeout_secs": 10 @@ -265,8 +283,6 @@ You can also pass `--help` to benchmarks to print the available options. | `--max-gpu-temp` | Maximum GPU temperature in °C. Too low slows benchmarks; too high increases noise. (default: 60) | | `--max-warming-secs` | Maximum seconds allowed for GPU warming before an error is thrown. (default: 60) | | `--max-cooling-secs` | Maximum seconds allowed for GPU cooling before an error is thrown. (default: 60) | -| `--output-hip-device-properties-context` | Output a `hip_device_properties` object in the context object, containing details about the GPU. | -| `--output-amdsmi-context` | Output an `amdsmi` object in the context object, containing details about the GPU. | | `--output-batches` | Output a `batches` array for each specialization, containing per-batch details. | | `--spaces-per-indent` | Number of spaces per indentation level in JSON output. Set to 0 for no indentation. (default: 4) | | `--stream-blocking-timeout-secs` | 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: 10) | @@ -344,13 +360,15 @@ Before benchmarking a specialization, primbench: primbench currently does not cool down the GPU *while* benchmarking a specialization, but this may be changed in the future. -Temperatures are read via AMD SMI. Warming uses short GPU workloads; cooling waits until the GPU naturally drops back into range. If either process takes more than 60 seconds, primbench aborts (`--max-warming-secs`, `--max-cooling-secs`). +Temperatures are read using AMD SMI/NVML. Warming uses short GPU workloads; cooling waits until the GPU naturally drops back into range. If either process takes more than 60 seconds, primbench aborts (`--max-warming-secs`, `--max-cooling-secs`). ### GPU Temperature Sensor Selection -primbench supports multiple GPU temperature sensors exposed by AMD SMI, such as *edge* and *hotspot*. At startup, it probes the available sensors and selects the first one that successfully returns a reading. The chosen sensor type is cached and used consistently for warming, cooling, and reporting. If no supported sensor can be read, primbench terminates with an error. +primbench supports multiple GPU temperature sensors exposed by AMD SMI, such as `edge` and `hotspot`. At startup, it probes the available sensors and selects the first one that successfully returns a reading. The chosen sensor type is cached and used consistently for warming, cooling, and reporting. If no supported sensor can be read, primbench terminates with an error. + +The selected temperature sensor type is recorded in the JSON output as `context.general.temperature_type`. -The selected temperature sensor type is recorded in the JSON output as `context.general.temp_type`. +primbench always uses the GPU temperature sensor called `gpu` for CUDA benchmarks. ### GPU Cache Clearing @@ -410,7 +428,7 @@ The command below embeds the current branch and commit hash. If the repository is in a detached HEAD state (for example in CI or when checking out a specific commit), no branch is active and `BRANCH_NAME` is set to `DETACHED`. ```bash -hipcc -o copy_benchmark examples/copy_benchmark.cpp \ +hipcc -o copy_benchmark examples/hip/copy_benchmark.cpp \ -I. \ -lamd_smi \ -DBRANCH_NAME=\"$(git symbolic-ref -q --short HEAD || echo DETACHED)\" \ diff --git a/shared/primbench/examples/cuda/copy_benchmark.cu b/shared/primbench/examples/cuda/copy_benchmark.cu new file mode 100644 index 000000000000..b2dd51729ef5 --- /dev/null +++ b/shared/primbench/examples/cuda/copy_benchmark.cu @@ -0,0 +1,98 @@ +#include "primbench.hpp" + +#include + +// All benchmarked types must be declared +// This allows you to for example format `long long` as `int64_t`, or `i64` +PRIMBENCH_REGISTER_TYPE(char, "char") +PRIMBENCH_REGISTER_TYPE(long long, "long long") + +// Simple copy kernel +template +__global__ __launch_bounds__(BlockSize) +void copy_kernel(const T* input, T* output) +{ + unsigned int idx = threadIdx.x + blockIdx.x * BlockSize * ItemsPerThread; +#pragma unroll + for(unsigned int i = 0; i < ItemsPerThread; ++i) + output[idx + i * BlockSize] = input[idx + i * BlockSize]; +} + +template +struct copy_benchmark : public primbench::benchmark_interface +{ + primbench::json meta() const override + { + // This meta() method must return an `algo` key + // The `algo` name must be identical across all queued specializations + // primbench::json objects can be nested + return primbench::json{}.add("algo", "copy").add("type", primbench::name()); + } + + // This contains both the setup and running of the benchmark + void run(primbench::state& state) override + { + const auto& stream = state.stream; + const auto& bytes = state.size; + + // primbench::log() calls report progress in gray + // They help with discovering slow setup steps + primbench::log("Calculating items"); + size_t N = bytes / sizeof(T); + constexpr unsigned int BlockSize = 256; + constexpr unsigned int ItemsPerThread = 4; + + const size_t items_per_block = BlockSize * ItemsPerThread; + const size_t items = items_per_block * ((N + items_per_block - 1) / items_per_block); + + primbench::log("Generating input data"); + std::vector h_input(items); + for(size_t i = 0; i < items; ++i) + h_input[i] = T(i); + + primbench::log("Allocating output vector"); + std::vector h_output(items); + + primbench::log("Allocating device memory"); + T* d_input; + T* d_output; + PRIMBENCH_CHECK(cudaMalloc(&d_input, items * sizeof(T))); + PRIMBENCH_CHECK(cudaMalloc(&d_output, items * sizeof(T))); + + primbench::log("Copying to device"); + PRIMBENCH_CHECK(cudaMemcpyAsync(d_input, + h_input.data(), + items * sizeof(T), + cudaMemcpyHostToDevice, + stream)); + + dim3 grid(items / items_per_block); + dim3 block(BlockSize); + + // primbench uses this to calculates the items/sec and bytes/sec + state.set_items(items); + state.add_reads(items); + state.add_writes(items); + + // This passes a lambda to primbench, which calls it many times + // primbench completely handles synchronization + state.run( + [&] { + copy_kernel + <<>>(d_input, d_output); + }); + + PRIMBENCH_CHECK(cudaFree(d_input)); + PRIMBENCH_CHECK(cudaFree(d_output)); + } +}; + +int main(int argc, char* argv[]) +{ + primbench::executor executor(argc, argv); + + executor.queue>(); + executor.queue>(); + + executor.run(); +} diff --git a/shared/primbench/examples/copy_benchmark.cpp b/shared/primbench/examples/hip/copy_benchmark.cpp similarity index 84% rename from shared/primbench/examples/copy_benchmark.cpp rename to shared/primbench/examples/hip/copy_benchmark.cpp index f8f2689bfea6..f926f205709a 100644 --- a/shared/primbench/examples/copy_benchmark.cpp +++ b/shared/primbench/examples/hip/copy_benchmark.cpp @@ -54,15 +54,15 @@ struct copy_benchmark : public primbench::benchmark_interface primbench::log("Allocating device memory"); T* d_input; T* d_output; - PRIMBENCH_HIP_CHECK(hipMalloc(&d_input, items * sizeof(T))); - PRIMBENCH_HIP_CHECK(hipMalloc(&d_output, items * sizeof(T))); + PRIMBENCH_CHECK(hipMalloc(&d_input, items * sizeof(T))); + PRIMBENCH_CHECK(hipMalloc(&d_output, items * sizeof(T))); primbench::log("Copying to device"); - PRIMBENCH_HIP_CHECK(hipMemcpyAsync(d_input, - h_input.data(), - items * sizeof(T), - hipMemcpyHostToDevice, - stream)); + PRIMBENCH_CHECK(hipMemcpyAsync(d_input, + h_input.data(), + items * sizeof(T), + hipMemcpyHostToDevice, + stream)); dim3 grid(items / items_per_block); dim3 block(BlockSize); @@ -80,8 +80,8 @@ struct copy_benchmark : public primbench::benchmark_interface <<>>(d_input, d_output); }); - PRIMBENCH_HIP_CHECK(hipFree(d_input)); - PRIMBENCH_HIP_CHECK(hipFree(d_output)); + PRIMBENCH_CHECK(hipFree(d_input)); + PRIMBENCH_CHECK(hipFree(d_output)); } }; diff --git a/shared/primbench/primbench.hpp b/shared/primbench/primbench.hpp index ad7093e4cf2b..cd11db316c5e 100644 --- a/shared/primbench/primbench.hpp +++ b/shared/primbench/primbench.hpp @@ -20,8 +20,16 @@ #pragma once -#include -#include +#ifdef __HIP__ + #include + #include +#elif defined(__CUDACC__) + #include + #include +#else + #error "Unsupported GPU platform" +#endif + #include #include @@ -31,7 +39,6 @@ #include #include #include -#include #include #include #include @@ -39,20 +46,16 @@ #include #include -/** - * \brief Default GPU cache size used for clearing caches. - * - * This conservative size is currently used to evict cached data before - * kernel launches. In the future, introducing HSA as a dependency may - * allow querying the actual largest GPU cache at runtime. - */ +/// Default GPU cache size used for clearing caches. +/// +/// This conservative size is currently used to evict cached data before +/// kernel launches. In the future, introducing HSA as a dependency may +/// allow querying the actual largest GPU cache at runtime. #ifndef PRIMBENCH_GPU_CACHE_SIZE #define PRIMBENCH_GPU_CACHE_SIZE 256 * primbench::MiB #endif -/** - * \brief Registers a custom type name, used by `primbench::name()`. - */ +/// Registers a custom type name, used by `primbench::name()`. #define PRIMBENCH_REGISTER_TYPE(TYPE, NAME) \ namespace primbench::detail \ { \ @@ -63,34 +66,60 @@ }; \ } -/** - * \brief Exits the program with an error message if the given HIP API call returns a failure - * status. - */ -#define PRIMBENCH_HIP_CHECK(status) \ - { \ - if(status != hipSuccess) \ - { \ - std::cerr << __FILE__ << ":" << __LINE__ \ - << ": HIP error: " << hipGetErrorString(status) << "\n"; \ - exit(status); \ - } \ - } - -/** - * \brief Exits the program with an error message if the given AMD SMI API call returns a - * failure status. - */ -#define PRIMBENCH_AMDSMI_CHECK(status) \ - { \ - if(status != AMDSMI_STATUS_SUCCESS) \ - { \ - const char* errstr = "(unknown)"; \ - amdsmi_status_code_to_string(status, &errstr); \ - std::cerr << __FILE__ << ":" << __LINE__ << ": AMDSMI error: " << errstr << "\n"; \ - std::exit(status); \ - } \ - } +#ifdef __HIP__ + /// Exits the program with an error message if the given API call returns a failure status. + #define PRIMBENCH_CHECK(status) \ + do \ + { \ + if(status != hipSuccess) \ + { \ + std::cerr << __FILE__ << ":" << __LINE__ \ + << ": HIP error: " << hipGetErrorString(status) << "\n"; \ + exit(status); \ + } \ + } \ + while(0) + + /// Exits the program with an error message if the given AMD SMI API call returns a failure status. + #define PRIMBENCH_AMDSMI_CHECK(status) \ + do \ + { \ + if(status != AMDSMI_STATUS_SUCCESS) \ + { \ + const char* errstr = "(unknown)"; \ + amdsmi_status_code_to_string(status, &errstr); \ + std::cerr << __FILE__ << ":" << __LINE__ << ": AMDSMI error: " << errstr << "\n"; \ + exit(status); \ + } \ + } \ + while(0) +#elif defined(__CUDACC__) + /// Exits the program with an error message if the given API call returns a failure status. + #define PRIMBENCH_CHECK(status) \ + do \ + { \ + if(status != cudaSuccess) \ + { \ + std::cerr << __FILE__ << ":" << __LINE__ \ + << ": CUDA error: " << cudaGetErrorString(status) << "\n"; \ + exit(status); \ + } \ + } \ + while(0) + + /// Exits the program with an error message if the given NVML API call returns a failure status. + #define PRIMBENCH_NVML_CHECK(status) \ + do \ + { \ + if(status != NVML_SUCCESS) \ + { \ + std::cerr << __FILE__ << ":" << __LINE__ \ + << ": NVML error: " << nvmlErrorString(status) << "\n"; \ + exit(status); \ + } \ + } \ + while(0) +#endif namespace primbench { @@ -99,52 +128,53 @@ inline constexpr size_t KiB = 1024; inline constexpr size_t MiB = 1024 * KiB; inline constexpr size_t GiB = 1024 * MiB; +// Backend-agnostic stream type and default stream constant. +#ifdef __HIP__ +using stream_t = hipStream_t; +constexpr stream_t default_stream = hipStreamDefault; +#elif defined(__CUDACC__) +using stream_t = cudaStream_t; +constexpr stream_t default_stream = cudaStreamDefault; +#endif + // Forward declarations for the detail namespace. extern const size_t MiB; template void log(Args&&... args); -/** - * \brief Settings that the user can change by passing arguments via their CLI. - */ +/// Settings that benchmarks and users can pass. struct settings { - size_t size = 128 * primbench::MiB; /**< Input array size */ - bool hot = false; /**< Hot means not clearing GPU cache between batches */ - uint32_t seed = 42; /**< The seed to use for input array generation */ - std::string json_out = "results.json"; /**< Output JSON file path */ - std::string csv_out = ""; /**< Output CSV file path */ - std::string filter = ""; /**< Regex filter of specialization names to benchmark */ - bool dry = false; /**< Flag to perform a dry run */ - double min_gpu_ms_per_batch = 10.0; /**< Minimum GPU batch duration */ - double min_secs = 1.0; /**< Minimum benchmark duration */ - double noise_timeout_secs = 10.0; /**< Max duration before noisy benchmark times out */ - size_t batch_window_size = 10; /**< Noise window size for early stopping */ - double noise_tolerance_percent = 1.0; /**< Noise tolerance for early stopping */ - uint16_t min_gpu_temp = 50; /**< Minimum GPU temperature */ - uint16_t max_gpu_temp = 60; /**< Maximum GPU temperature */ - double max_warming_secs = 60.0; /**< Max GPU warmup time */ - double max_cooling_secs = 60.0; /**< Max GPU cooldown time */ - bool output_hip_device_properties_context - = false; /**< Flag to output HIP device properties context */ - bool output_amdsmi_context = false; /**< Flag to output AMD SMI context */ - bool output_batches = false; /**< Flag to output batch details */ - uint32_t spaces_per_indent = 4; /**< JSON indentation spaces */ - double stream_blocking_timeout_secs - = 10.0; /**< Max duration before stream blocking times out */ + size_t size = 128 * primbench::MiB; ///< Input array size. + bool hot = false; ///< Hot means not clearing GPU cache between batches. + uint32_t seed = 42; ///< The seed to use for input array generation. + std::string json_out = "results.json"; ///< Output JSON file path. + std::string csv_out = ""; ///< Output CSV file path. + std::string filter = ""; ///< Regex filter of specialization names to benchmark. + bool dry = false; ///< Flag to perform a dry run. + double min_gpu_ms_per_batch = 10.0; ///< Minimum GPU batch duration. + double min_secs = 1.0; ///< Minimum benchmark duration. + double noise_timeout_secs = 10.0; ///< Max duration before noisy benchmark times out. + size_t batch_window_size = 10; ///< Noise window size for early stopping. + double noise_tolerance_percent = 1.0; ///< Noise tolerance for early stopping. + uint16_t min_gpu_temp = 50; ///< Minimum GPU temperature. + uint16_t max_gpu_temp = 60; ///< Maximum GPU temperature. + double max_warming_secs = 60.0; ///< Max GPU warmup time. + double max_cooling_secs = 60.0; ///< Max GPU cooldown time. + bool output_batches = false; ///< Flag to output batch details. + uint32_t spaces_per_indent = 4; ///< JSON indentation spaces. + double stream_blocking_timeout_secs = 10.0; ///< Max duration before stream blocking times out. using custom_arg_value = std::variant; std::map - custom_args; /**< Custom user-registered arguments with types */ + custom_args; ///< Custom user-registered arguments with types. }; // struct settings namespace detail { -/** - * \brief Fallback for primbench::name(). - */ +/// Fallback for primbench::name(). template struct type_name { @@ -153,11 +183,8 @@ struct type_name static inline const char* name = ""; }; -/** - * \brief Caches whether ANSI color output is enabled. - * - * The standard is described at https://bixense.com/clicolors/ - */ +/// Caches whether ANSI color output is enabled. +/// The standard is described at https://bixense.com/clicolors/ inline bool use_color() { static const bool result = [] @@ -171,9 +198,7 @@ inline bool use_color() return result; } -/** - * \brief Clears the current line if colors are enabled, otherwise prints a newline. - */ +/// Clears the current line if colors are enabled, otherwise prints a newline. inline std::ostream& clearline(std::ostream& os) { if(use_color()) @@ -183,9 +208,7 @@ inline std::ostream& clearline(std::ostream& os) return os; } -/** - * \brief Resets the output text formatting to default if colors are enabled. - */ +/// Resets the output text formatting to default if colors are enabled. inline std::ostream& reset(std::ostream& os) { if(use_color()) @@ -193,9 +216,7 @@ inline std::ostream& reset(std::ostream& os) return os; } -/** - * \brief Sets the output text color to gray if colors are enabled. - */ +/// Sets the output text color to gray if colors are enabled. inline std::ostream& gray(std::ostream& os) { if(use_color()) @@ -203,9 +224,7 @@ inline std::ostream& gray(std::ostream& os) return os; } -/** - * \brief Sets the output text color to green if colors are enabled. - */ +/// Sets the output text color to green if colors are enabled. inline std::ostream& green(std::ostream& os) { if(use_color()) @@ -213,9 +232,7 @@ inline std::ostream& green(std::ostream& os) return os; } -/** - * \brief Sets the output text color to red if colors are enabled. - */ +/// Sets the output text color to red if colors are enabled. inline std::ostream& red(std::ostream& os) { if(use_color()) @@ -223,9 +240,7 @@ inline std::ostream& red(std::ostream& os) return os; } -/** - * \brief Sets the output text color to yellow if colors are enabled. - */ +/// Sets the output text color to yellow if colors are enabled. inline std::ostream& yellow(std::ostream& os) { if(use_color()) @@ -233,9 +248,7 @@ inline std::ostream& yellow(std::ostream& os) return os; } -/** - * \brief Sets the output text color to blue if colors are enabled. - */ +/// Sets the output text color to blue if colors are enabled. inline std::ostream& blue(std::ostream& os) { if(use_color()) @@ -243,200 +256,236 @@ inline std::ostream& blue(std::ostream& os) return os; } -/** - * \brief Wrapper for AMD SMI (System Management Interface) GPU monitoring. - * - * Initializes AMD GPU metrics, clocks, and memory usage on construction, - * and shuts down AMD SMI on destruction. Provides methods to: - * - Retrieve current GPU statistics (`get_stats`) - * - Serialize GPU statistics and device context to JSON (`serialize_stats`, `serialize_context`) - * - Read GPU temperature (`get_temp`) - * - * Contains internal types for holding GPU stats and context information. - */ -class amdsmi +/// Backend-agnostic wrappers for HIP and CUDA. +/// The functions in this namespace are sorted alphabetically. +namespace gpu_backend +{ + +#ifdef __HIP__ +using error_t = hipError_t; +using event_t = hipEvent_t; +constexpr unsigned host_register_mapped = hipHostRegisterMapped; +#elif defined(__CUDACC__) +using error_t = cudaError_t; +using event_t = cudaEvent_t; +constexpr unsigned host_register_mapped = cudaHostRegisterMapped; +#endif + +#ifdef __HIP__ + #define PRIMBENCH_PREFIX_BACKEND(fn) hip##fn +#elif defined(__CUDACC__) + #define PRIMBENCH_PREFIX_BACKEND(fn) cuda##fn +#endif + +inline error_t event_create(event_t* event) +{ + return PRIMBENCH_PREFIX_BACKEND(EventCreate)(event); +} + +inline error_t event_destroy(event_t event) +{ + return PRIMBENCH_PREFIX_BACKEND(EventDestroy)(event); +} + +inline error_t event_elapsed_time(float* ms, event_t start, event_t stop) +{ + return PRIMBENCH_PREFIX_BACKEND(EventElapsedTime)(ms, start, stop); +} + +inline error_t event_record(event_t event, stream_t stream) +{ + return PRIMBENCH_PREFIX_BACKEND(EventRecord)(event, stream); +} + +inline error_t get_last_error() +{ + return PRIMBENCH_PREFIX_BACKEND(GetLastError)(); +} + +inline error_t gpu_free(void* device) +{ + return PRIMBENCH_PREFIX_BACKEND(Free)(device); +} + +inline error_t gpu_malloc(void** device, size_t size) +{ + return PRIMBENCH_PREFIX_BACKEND(Malloc)(device, size); +} + +inline error_t host_get_device_pointer(void** device, void* host, unsigned flags) +{ + return PRIMBENCH_PREFIX_BACKEND(HostGetDevicePointer)(device, host, flags); +} + +inline error_t host_register(void* ptr, size_t size, unsigned flags) +{ + return PRIMBENCH_PREFIX_BACKEND(HostRegister)(ptr, size, flags); +} + +inline error_t host_unregister(void* ptr) +{ + return PRIMBENCH_PREFIX_BACKEND(HostUnregister)(ptr); +} + +inline error_t memset_async(void* device, int value, size_t count, stream_t stream) +{ + return PRIMBENCH_PREFIX_BACKEND(MemsetAsync)(device, value, count, stream); +} + +inline error_t stream_create(stream_t* stream) +{ + return PRIMBENCH_PREFIX_BACKEND(StreamCreate)(stream); +} + +inline error_t stream_destroy(stream_t stream) +{ + return PRIMBENCH_PREFIX_BACKEND(StreamDestroy)(stream); +} + +inline error_t stream_synchronize(stream_t stream) +{ + return PRIMBENCH_PREFIX_BACKEND(StreamSynchronize)(stream); +} + +#undef PRIMBENCH_PREFIX_BACKEND + +} // namespace gpu_backend + +/// Bring GPU backend wrappers into this namespace so we can call +/// functions like gpu_malloc() without the gpu_backend:: prefix. +using namespace gpu_backend; + +#ifdef __HIP__ + +/// Wrapper for HIP and AMD SMI (System Management Interface) GPU monitoring. +class monitor { public: - // Delete all copy/move constructors and assignment operators - amdsmi(const amdsmi&) = delete; - amdsmi& operator=(const amdsmi&) = delete; - amdsmi(amdsmi&&) = delete; - amdsmi& operator=(amdsmi&&) = delete; - - // Singleton accessor - static amdsmi& instance() + // Delete all copy/move constructors and assignment operators. + monitor(const monitor&) = delete; + monitor& operator=(const monitor&) = delete; + monitor(monitor&&) = delete; + monitor& operator=(monitor&&) = delete; + + /// Singleton accessor. + static monitor& instance() { - static amdsmi instance; + static monitor instance; return instance; } - /** - * \brief Represents GPU statistics including metrics, clocks, and VRAM usage. - */ - struct stats + uint16_t get_temp() const { - amdsmi_gpu_metrics_t metrics; - - // Clocks (current frequencies in MHz) - std::unordered_map> clocks; - - std::optional vram_used_bytes; - }; + int64_t t; + PRIMBENCH_AMDSMI_CHECK(amdsmi_get_temp_metric(m_amdsmi_device, + get_temperature_type(), + AMDSMI_TEMP_CURRENT, + &t)); + return t; + } - /** - * \brief Converts a stats object to a JSON string. - * \param stats The stats object to serialize. - * \return JSON string representing the stats. - */ - std::string serialize_stats(const stats& stats) const + std::string serialize_gpu_info() const { std::ostringstream ss; ss << "{"; - ss << "\"vram_used_bytes\":" << serialize_optional(stats.vram_used_bytes); + hipDeviceProp_t dev_prop; + PRIMBENCH_CHECK(hipGetDeviceProperties(&dev_prop, m_hip_device)); + ss << "\"name\":\"" << dev_prop.name << "\""; - if(has_clock(stats.clocks)) + std::string arch = dev_prop.gcnArchName; + size_t pos = arch.find(':'); + if(pos != std::string::npos) { - ss << ",\"clocks_mhz\":{"; - bool first = true; - for(const auto& kv : stats.clocks) - { - if(!first) - ss << ","; - ss << "\"" << kv.first << "\":" << serialize_optional(kv.second); - first = false; - } - ss << "}"; + arch = arch.substr(0, pos); } + ss << ",\"arch\":\"" << arch << "\""; - ss << ",\"metrics\":" << serialize_metrics(stats.metrics); + char pci_bus_id_arr[32]; + PRIMBENCH_CHECK(hipDeviceGetPCIBusId(pci_bus_id_arr, sizeof(pci_bus_id_arr), m_hip_device)); + ss << ",\"pci_bus_id\":\"" << pci_bus_id_arr << "\""; ss << "}"; return ss.str(); } - /** - * \brief Converts the GPU context to a JSON string. - * \return JSON string representing the context and current GPU state. - */ - std::string serialize_context() const + std::string serialize_backend_info() const { - const auto& ctx = m_context; - std::ostringstream ss; ss << "{"; - ss << "\"identity\":{"; - - ss << "\"product_name\":" << serialize_optional(ctx.product_name); - - ss << ",\"version\":" << serialize_optional(ctx.amdsmi_version); - - ss << ",\"metrics_version\":{"; - ss << "\"format\":" << std::to_string(ctx.amdsmi_metrics_version.format_revision); - ss << ",\"content\":" << std::to_string(ctx.amdsmi_metrics_version.content_revision); - ss << "}"; + // Backend name. + ss << "\"name\":\"hip\""; - ss << "}"; // End of "identity" object. + // HIP version. + ss << ",\"hip_version\":\"" << HIP_VERSION_MAJOR << "." << HIP_VERSION_MINOR << "." + << HIP_VERSION_PATCH << "-" << HIP_VERSION_GITHASH << "\""; - ss << ",\"power_cap\":{"; - ss << "\"current_microwatts\":" << serialize_optional(ctx.power_cap); - ss << ",\"default_microwatts\":" << serialize_optional(ctx.power_cap_default); - ss << ",\"dpm_mhz\":" << serialize_optional(ctx.power_cap_dpm); + // HIP runtime version (integer, e.g., 60443482 -> 6.4.43482). + int runtime_ver; + PRIMBENCH_CHECK(hipRuntimeGetVersion(&runtime_ver)); + int major = runtime_ver / 10000000; + int minor = (runtime_ver / 100000) % 100; + int patch = runtime_ver % 100000; + ss << ",\"runtime_version\":\"" << major << "." << minor << "." << patch << "\""; + + // HIP driver version (integer, same format). + int driver_ver; + PRIMBENCH_CHECK(hipDriverGetVersion(&driver_ver)); + major = driver_ver / 10000000; + minor = (driver_ver / 100000) % 100; + patch = driver_ver % 100000; + ss << ",\"driver_version\":\"" << major << "." << minor << "." << patch << "\""; + + // Compiler. + ss << ",\"compiler\":{"; + ss << "\"name\":\"clang\""; + ss << ",\"version\":\"" << __clang_version__ << "\""; ss << "}"; - ss << ",\"vram\":{"; - ss << "\"vendor\":" << serialize_optional(ctx.vram_vendor); - ss << ",\"total_bytes\":" << serialize_optional(ctx.vram_total_bytes); + // Monitoring. + ss << ",\"monitoring\":{"; + ss << "\"name\":\"amdsmi\""; + amdsmi_version_t amdsmi_version; + PRIMBENCH_AMDSMI_CHECK(amdsmi_get_lib_version(&amdsmi_version)); + ss << ",\"version\":\"" << amdsmi_version.build << "\""; ss << "}"; - if(has_clock(ctx.clocks)) - { - ss << ",\"clocks\":{"; - bool first = true; - for(const auto& kv : ctx.clocks) - { - if(!first) - ss << ","; - ss << "\"" << kv.first << "\":"; - if(kv.second) - ss << "{" << "\"min_mhz\":" << kv.second->first - << ",\"max_mhz\":" << kv.second->second << "}"; - else - ss << "null"; - first = false; - } - ss << "}"; - } - - ss << ",\"stats\":" << serialize_stats(ctx.stats); - ss << "}"; return ss.str(); } - /** - * \brief Retrieve current GPU statistics. - * \return stats object containing current metrics, clocks, and memory usage. - */ - stats get_stats() const + std::string get_used_temperature_type_name() const { - stats stats{}; - - // Copy all GPU metrics - amdsmi_gpu_metrics_t metrics{}; - if(amdsmi_get_gpu_metrics_info(m_target, &metrics) == AMDSMI_STATUS_SUCCESS) - stats.metrics = metrics; + return get_temperature_type_name(get_temperature_type()); + } - // Clocks - for(auto clk : clk_types) - { - amdsmi_clk_info_t clk_info{}; - if(amdsmi_get_clock_info(m_target, clk, &clk_info) == AMDSMI_STATUS_SUCCESS) - stats.clocks[clk_type_to_string(clk)] = clk_info.clk; - } +private: + monitor() + { + PRIMBENCH_CHECK(hipGetDevice(&m_hip_device)); - // Memory usage - uint64_t vram_used; - if(amdsmi_get_gpu_memory_usage(m_target, AMDSMI_MEM_TYPE_VRAM, &vram_used) - == AMDSMI_STATUS_SUCCESS) - stats.vram_used_bytes = vram_used; + PRIMBENCH_AMDSMI_CHECK(amdsmi_init(AMDSMI_INIT_AMD_GPUS)); - return stats; + // This can't be run in a member initializer list, + // since the above amdsmi_init() call must be run first. + m_amdsmi_device = get_amdsmi_device(); } - /** - * \brief Converts a temperature type enum to a string. - * \param type Temperature type enum value. - * \return String representation of the temperature type. - * \note Exits with failure if the temperature type is not recognized. - */ - static const char* get_temp_type_name(amdsmi_temperature_type_t type) + ~monitor() { - switch(type) - { - case AMDSMI_TEMPERATURE_TYPE_EDGE: return "edge"; - case AMDSMI_TEMPERATURE_TYPE_HOTSPOT: return "hotspot"; - default: - std::cerr << "\nError: Failed to match temperature type " << type - << " to a string\n"; - exit(EXIT_FAILURE); - } + PRIMBENCH_AMDSMI_CHECK(amdsmi_shut_down()); } - /** - * \brief Determines the first available GPU temperature sensor type. - * - * Attempts to read temperature from multiple sensor types in priority order. - * Returns the first type that successfully provides a temperature reading. - * - * \return The first successfully queried temperature sensor type. - * \note Exits with failure if no temperature sensors are accessible. - * The result is cached as a static variable. - */ - amdsmi_temperature_type_t get_temp_type() const + /// Determines the first available GPU temperature sensor type. + /// + /// Attempts to read temperature from multiple sensor types in priority order. + /// Returns the first type that successfully provides a temperature reading. + /// + /// The result is cached in a static variable. + amdsmi_temperature_type_t get_temperature_type() const { - static const amdsmi_temperature_type_t temp_type = [&] + static const amdsmi_temperature_type_t temperature_type = [&] { const amdsmi_temperature_type_t types[] = { AMDSMI_TEMPERATURE_TYPE_EDGE, @@ -446,7 +495,7 @@ class amdsmi int64_t t; for(auto type : types) { - if(amdsmi_get_temp_metric(m_target, type, AMDSMI_TEMP_CURRENT, &t) + if(amdsmi_get_temp_metric(m_amdsmi_device, type, AMDSMI_TEMP_CURRENT, &t) == AMDSMI_STATUS_SUCCESS) { return type; @@ -460,477 +509,192 @@ class amdsmi if(!first) std::cerr << ", "; first = false; - std::cerr << get_temp_type_name(type); + std::cerr << get_temperature_type_name(type); } std::cerr << "\n"; exit(EXIT_FAILURE); }(); - return temp_type; - } - - /** - * \brief Reads the GPU temperature. - * \return Temperature in °C. - */ - uint16_t get_temp() const - { - int64_t t = 0; - PRIMBENCH_AMDSMI_CHECK( - amdsmi_get_temp_metric(m_target, get_temp_type(), AMDSMI_TEMP_CURRENT, &t)); - return t; + return temperature_type; } -private: - amdsmi() + /// Converts a temperature type enum to a string. + static const char* get_temperature_type_name(amdsmi_temperature_type_t type) { - PRIMBENCH_AMDSMI_CHECK(amdsmi_init(AMDSMI_INIT_AMD_GPUS)); - - // These can't be turned into a member initializer list, - // because the above amdsmi_init() call must happen first. - m_target = get_target(); - m_context = get_context(m_target); - } - - ~amdsmi() - { - PRIMBENCH_AMDSMI_CHECK(amdsmi_shut_down()); - } - - /** - * \brief Holds detailed information about the GPU and AMD SMI context. - */ - struct context - { - std::optional product_name; - - std::string amdsmi_version; - - struct + switch(type) { - uint8_t format_revision; - uint8_t content_revision; - } amdsmi_metrics_version; - - // Clocks (min/max in MHz) - std::unordered_map>> clocks; - - std::optional power_cap; - std::optional power_cap_default; - std::optional power_cap_dpm; - - std::optional vram_vendor; - std::optional vram_total_bytes; - - stats stats; - }; - - /** - * \brief Clock types queried for the GPU. - * \note Used internally to query all supported GPU clock domains (sys, mem, soc, etc.). - */ - const std::vector clk_types = { - AMDSMI_CLK_TYPE_SYS, - AMDSMI_CLK_TYPE_DF, - AMDSMI_CLK_TYPE_DCEF, - AMDSMI_CLK_TYPE_SOC, - AMDSMI_CLK_TYPE_MEM, - AMDSMI_CLK_TYPE_PCIE, - AMDSMI_CLK_TYPE_VCLK0, - AMDSMI_CLK_TYPE_VCLK1, - AMDSMI_CLK_TYPE_DCLK0, - AMDSMI_CLK_TYPE_DCLK1, - }; + case AMDSMI_TEMPERATURE_TYPE_EDGE: return "edge"; + case AMDSMI_TEMPERATURE_TYPE_HOTSPOT: return "hotspot"; + default: + std::cerr << "\nError: Failed to match temperature type " << type + << " to a string\n"; + exit(EXIT_FAILURE); + } + } - /** - * \brief Finds the AMD SMI processor handle matching the current HIP device. - * \return AMD SMI processor handle of the target GPU. - */ - amdsmi_processor_handle get_target() const + /// Finds the AMD SMI processor handle matching the current HIP device. + amdsmi_processor_handle get_amdsmi_device() const { - int hip_dev; - PRIMBENCH_HIP_CHECK(hipGetDevice(&hip_dev)); - hipDeviceProp_t hip_props; - PRIMBENCH_HIP_CHECK(hipGetDeviceProperties(&hip_props, hip_dev)); + PRIMBENCH_CHECK(hipGetDeviceProperties(&hip_props, m_hip_device)); - // Build the AMD SMI BDF struct from HIP device properties + // Build the AMD SMI BDF struct from HIP device properties. amdsmi_bdf_t addr{ - .function_number = 0, // HIP doesn't expose PCI function ID + .function_number = 0, // HIP doesn't expose PCI function ID. .device_number = static_cast(hip_props.pciDeviceID), .bus_number = static_cast(hip_props.pciBusID), .domain_number = static_cast(hip_props.pciDomainID), }; - amdsmi_processor_handle target; - PRIMBENCH_AMDSMI_CHECK(amdsmi_get_processor_handle_from_bdf(addr, &target)); + amdsmi_processor_handle amdsmi_device; + PRIMBENCH_AMDSMI_CHECK(amdsmi_get_processor_handle_from_bdf(addr, &amdsmi_device)); - return target; + return amdsmi_device; } - /** - * \brief Builds a context object with GPU details and metrics. - * \param target AMD SMI processor handle of the target GPU. - * \return Context object with detailed GPU information. - */ - context get_context(amdsmi_processor_handle target) const - { - context ctx{}; - - amdsmi_board_info_t board_info; - if(amdsmi_get_gpu_board_info(target, &board_info) == AMDSMI_STATUS_SUCCESS) - ctx.product_name = board_info.product_name; - - amdsmi_version_t amdsmi_version; - if(amdsmi_get_lib_version(&amdsmi_version) == AMDSMI_STATUS_SUCCESS) - ctx.amdsmi_version = amdsmi_version.build; - - amdsmi_gpu_metrics_t metrics{}; - if(amdsmi_get_gpu_metrics_info(target, &metrics) == AMDSMI_STATUS_SUCCESS) - { - ctx.amdsmi_metrics_version.format_revision = metrics.common_header.format_revision; - ctx.amdsmi_metrics_version.content_revision = metrics.common_header.content_revision; - } - - for(auto clk : clk_types) - { - amdsmi_clk_info_t clk_info{}; - if(amdsmi_get_clock_info(target, clk, &clk_info) == AMDSMI_STATUS_SUCCESS) - ctx.clocks[clk_type_to_string(clk)] - = std::make_pair(clk_info.min_clk, clk_info.max_clk); - } - - amdsmi_power_cap_info_t pcap; - if(amdsmi_get_power_cap_info(target, 0, &pcap) == AMDSMI_STATUS_SUCCESS) - { - ctx.power_cap = pcap.power_cap; - ctx.power_cap_default = pcap.default_power_cap; - ctx.power_cap_dpm = pcap.dpm_cap; - } + int m_hip_device; ///< HIP device. + amdsmi_processor_handle m_amdsmi_device; ///< AMD SMI device. - char vram_vendor_buf[128]; - if(amdsmi_get_gpu_vram_vendor(target, vram_vendor_buf, sizeof(vram_vendor_buf)) - == AMDSMI_STATUS_SUCCESS) - ctx.vram_vendor = vram_vendor_buf; +}; // class monitor - uint64_t vram_total; - if(amdsmi_get_gpu_memory_total(target, AMDSMI_MEM_TYPE_VRAM, &vram_total) - == AMDSMI_STATUS_SUCCESS) - ctx.vram_total_bytes = vram_total; +#elif defined(__CUDACC__) - ctx.stats = get_stats(); - - return ctx; - } +/// Wrapper for CUDA and NVML (NVIDIA Management Library) GPU monitoring. +class monitor +{ +public: + // Delete all copy/move constructors and assignment operators. + monitor(const monitor&) = delete; + monitor& operator=(const monitor&) = delete; + monitor(monitor&&) = delete; + monitor& operator=(monitor&&) = delete; - /** - * \brief Converts a clock type enum to a string. - * \param clk Clock type. - * \return Corresponding string representation of the clock type. - */ - std::string clk_type_to_string(amdsmi_clk_type_t clk) const + /// Singleton accessor. + static monitor& instance() { - switch(clk) - { - case AMDSMI_CLK_TYPE_SYS: return "sys"; - case AMDSMI_CLK_TYPE_DF: return "df"; - case AMDSMI_CLK_TYPE_DCEF: return "dcef"; - case AMDSMI_CLK_TYPE_SOC: return "soc"; - case AMDSMI_CLK_TYPE_MEM: return "mem"; - case AMDSMI_CLK_TYPE_PCIE: return "pcie"; - case AMDSMI_CLK_TYPE_VCLK0: return "vclk0"; - case AMDSMI_CLK_TYPE_VCLK1: return "vclk1"; - case AMDSMI_CLK_TYPE_DCLK0: return "dclk0"; - case AMDSMI_CLK_TYPE_DCLK1: return "dclk1"; - } - std::cerr << "Error: Failed to match clock type " << clk << " to a string\n"; - exit(EXIT_FAILURE); + static monitor instance; + return instance; } - /** - * \brief Serializes an optional string to JSON format. - * \param opt Optional string to serialize. - * \return JSON string or "null" if empty. - */ - std::string serialize_optional(const std::optional& opt) const + uint16_t get_temp() const { - return opt ? ("\"" + *opt + "\"") : "null"; + unsigned int t; + PRIMBENCH_NVML_CHECK(nvmlDeviceGetTemperature(m_nvml_device, NVML_TEMPERATURE_GPU, &t)); + return t; } - /** - * \brief Serializes an optional numeric value to JSON format. - * \tparam T Numeric type. - * \param opt Optional value to serialize. - * \return JSON numeric string or "null" if empty. - */ - template - std::string serialize_optional(const std::optional& opt) const + std::string serialize_gpu_info() const { - return opt ? std::to_string(*opt) : "null"; - } + std::ostringstream ss; + ss << "{"; - /** - * \brief Checks if any clock frequency is available in the provided clock map. - * \tparam T The value type of the optional (either uint32_t or std::pair). - * \param clocks Map of clock names to optional frequency values. - * \return true if at least one clock has a value, false if all are std::nullopt. - */ - template - static bool has_clock(const std::unordered_map>& clocks) - { - return std::any_of(clocks.begin(), - clocks.end(), - [](const auto& kv) { return kv.second.has_value(); }); + cudaDeviceProp dev_prop; + PRIMBENCH_CHECK(cudaGetDeviceProperties(&dev_prop, m_cuda_device)); + ss << "\"name\":\"" << dev_prop.name << "\""; + + ss << ",\"arch\":\"" << dev_prop.major << "." << dev_prop.minor << "\""; + + char pci_bus_id_arr[32]; + PRIMBENCH_CHECK( + cudaDeviceGetPCIBusId(pci_bus_id_arr, sizeof(pci_bus_id_arr), m_cuda_device)); + ss << ",\"pci_bus_id\":\"" << pci_bus_id_arr << "\""; + + ss << "}"; + return ss.str(); } - /** - * \brief Serializes the raw GPU metrics to a JSON string. - * \param metrics GPU metrics to serialize. - * \return JSON string representing all available metrics. - * \note Serialization includes fields conditionally depending on the metrics content revision. - */ - std::string serialize_metrics(const amdsmi_gpu_metrics_t& metrics) const + std::string serialize_backend_info() const { std::ostringstream ss; ss << "{"; - bool first = true; - auto add_comma = [&]() - { - if(!first) - ss << ","; - first = false; - }; + // Backend name. + ss << "\"name\":\"cuda\""; + + // CUDA runtime version. + int major = CUDART_VERSION / 1000; + int minor = (CUDART_VERSION % 1000) / 10; + ss << ",\"runtime_version\":\"" << major << "." << minor << "\""; + + // CUDA driver version (integer, e.g., 9020 -> 9.2). + int driver_ver; + PRIMBENCH_CHECK(cudaDriverGetVersion(&driver_ver)); + major = driver_ver / 1000; + minor = (driver_ver % 1000) / 10; + ss << ",\"driver_version\":\"" << major << "." << minor << "\""; + + // Compiler. + ss << ",\"compiler\":{"; + ss << "\"name\":\"nvcc\""; + ss << ",\"version\":\"" << __CUDACC_VER_MAJOR__ << "." << __CUDACC_VER_MINOR__ << "." + << __CUDACC_VER_BUILD__ << "\""; + ss << "}"; - auto add_field = [&](const char* name, const auto& value) - { - add_comma(); - ss << "\"" << name << "\":" << value; - }; + // Monitoring. + ss << ",\"monitoring\":{"; + ss << "\"name\":\"nvml\""; + char nvml_ver[NVML_SYSTEM_NVML_VERSION_BUFFER_SIZE]; + PRIMBENCH_NVML_CHECK(nvmlSystemGetNVMLVersion(nvml_ver, sizeof(nvml_ver))); + ss << ",\"version\":\"" << nvml_ver << "\""; + ss << "}"; - auto add_array = [&](const char* name, auto&& arr) - { - add_comma(); - ss << "\"" << name << "\":["; - for(size_t i = 0; i < (sizeof(arr) / sizeof(*arr)); ++i) - { - if(i > 0) - ss << ","; - ss << arr[i]; - } - ss << "]"; - }; + ss << "}"; + return ss.str(); + } - struct revision_block - { - int min_content_revision; - std::function serialize; - }; + std::string get_used_temperature_type_name() const + { + return "gpu"; + } - std::vector blocks = { - {0, - [&] - { - add_field("average_socket_power_watts", metrics.average_socket_power); - add_field("energy_accumulator", metrics.energy_accumulator); - add_field("system_clock_counter_ns", metrics.system_clock_counter); - add_field("throttle_status", metrics.throttle_status); - add_field("current_fan_speed_rpm", metrics.current_fan_speed); - - ss << ",\"average_activity_percent\":{"; - first = true; - add_field("gfx", metrics.average_gfx_activity); - add_field("umc", metrics.average_umc_activity); - add_field("mm", metrics.average_mm_activity); - ss << "}"; - - ss << ",\"temp_celsius\":{"; - first = true; - add_field("edge", metrics.temperature_edge); - add_field("hotspot", metrics.temperature_hotspot); - add_field("mem", metrics.temperature_mem); - add_field("vrgfx", metrics.temperature_vrgfx); - add_field("vrsoc", metrics.temperature_vrsoc); - add_field("vrmem", metrics.temperature_vrmem); - ss << "}"; - - ss << ",\"average_frequency_mhz\":{"; - first = true; - add_field("gfxclk", metrics.average_gfxclk_frequency); - add_field("socclk", metrics.average_socclk_frequency); - add_field("uclk", metrics.average_uclk_frequency); - add_field("vclk0", metrics.average_vclk0_frequency); - add_field("dclk0", metrics.average_dclk0_frequency); - add_field("vclk1", metrics.average_vclk1_frequency); - add_field("dclk1", metrics.average_dclk1_frequency); - ss << "}"; - - ss << ",\"current_frequency_mhz\":{"; - first = true; - add_field("gfxclk", metrics.current_gfxclk); - add_field("socclk", metrics.current_socclk); - add_field("uclk", metrics.current_uclk); - add_field("vclk0", metrics.current_vclk0); - add_field("dclk0", metrics.current_dclk0); - add_field("vclk1", metrics.current_vclk1); - add_field("dclk1", metrics.current_dclk1); - ss << "}"; - - ss << ",\"pcie_link\":{"; - first = true; - add_field("pcie_link_width", metrics.pcie_link_width); - add_field("pcie_link_speed", metrics.pcie_link_speed); - ss << "}"; - } }, - {1, - [&] - { - add_field("gfx_activity_acc", metrics.gfx_activity_acc); - add_field("mem_activity_acc", metrics.mem_activity_acc); - add_array("hbm_temp_celsius", metrics.temperature_hbm); - } }, - {2, [&] { add_field("firmware_timestamp", metrics.firmware_timestamp); }}, - {3, - [&] - { - add_field("indep_throttle_status", metrics.indep_throttle_status); - - ss << ",\"voltage_mv\":{"; - first = true; - add_field("soc", metrics.voltage_soc); - add_field("gfx", metrics.voltage_gfx); - add_field("mem", metrics.voltage_mem); - ss << "}"; - } }, - {4, - [&] - { - add_field("current_socket_power", metrics.current_socket_power); - add_field("gfxclk_lock_status", metrics.gfxclk_lock_status); - - ss << ",\"xgmi_link\":{"; - first = true; - add_field("width", metrics.xgmi_link_width); - add_field("speed", metrics.xgmi_link_speed); - ss << "}"; - - ss << ",\"pcie_bandwidth\":{"; - first = true; - add_field("acc", metrics.pcie_bandwidth_acc); - add_field("inst", metrics.pcie_bandwidth_inst); - ss << "}"; - - ss << ",\"pcie_count_acc\":{"; - first = true; - add_field("l0_to_recov", metrics.pcie_l0_to_recov_count_acc); - add_field("replay", metrics.pcie_replay_count_acc); - add_field("replay_rover", metrics.pcie_replay_rover_count_acc); - ss << "}"; - - ss << ",\"xgmi_data_acc\":{"; - first = true; - add_array("read", metrics.xgmi_read_data_acc); - add_array("write", metrics.xgmi_write_data_acc); - ss << "}"; - - ss << ",\"current\":{"; - first = true; - add_array("gfxclks", metrics.current_gfxclks); - add_array("socclks", metrics.current_socclks); - add_array("vclk0s", metrics.current_vclk0s); - add_array("dclk0s", metrics.current_dclk0s); - ss << "}"; - - add_array("vcn_activity", metrics.vcn_activity); - } }, - {5, - [&] - { - add_array("jpeg_activity_percent", metrics.jpeg_activity); - - ss << ",\"pcie_nak_count_acc\":{"; - first = true; - add_field("sent", metrics.pcie_nak_sent_count_acc); - add_field("rcvd", metrics.pcie_nak_rcvd_count_acc); - ss << "}"; - } }, - {6, - [&] - { - add_field("accumulation_counter", metrics.accumulation_counter); - add_field("num_partition", metrics.num_partition); - add_field("pcie_lc_perf_other_end_recovery", - metrics.pcie_lc_perf_other_end_recovery); - - ss << ",\"residency_acc\":{"; - first = true; - add_field("prochot", metrics.prochot_residency_acc); - add_field("ppt", metrics.ppt_residency_acc); - add_field("socket_thm", metrics.socket_thm_residency_acc); - add_field("vr_thm", metrics.vr_thm_residency_acc); - add_field("hbm_thm", metrics.hbm_thm_residency_acc); - ss << "}"; - - // xcp_stats is too annoying and unimportant to serialize. - } }, - {7, - [&] - { - add_field("vram_max_bandwidth", metrics.vram_max_bandwidth); - add_array("xgmi_link_status", metrics.xgmi_link_status); - } }, - }; +private: + monitor() + { + PRIMBENCH_CHECK(cudaGetDevice(&m_cuda_device)); - int current_revision = m_context.amdsmi_metrics_version.content_revision; - for(auto& block : blocks) - { - if(current_revision < block.min_content_revision) - break; - block.serialize(); - } + PRIMBENCH_NVML_CHECK(nvmlInit()); - ss << "}"; - return ss.str(); + PRIMBENCH_NVML_CHECK(nvmlDeviceGetHandleByIndex(0, &m_nvml_device)); + } + + ~monitor() + { + PRIMBENCH_NVML_CHECK(nvmlShutdown()); } - amdsmi_processor_handle m_target; /**< AMD SMI handle of the GPU. */ - context m_context; /**< Cached GPU context and stats. */ + int m_cuda_device; ///< CUDA device. + nvmlDevice_t m_nvml_device; ///< NVML device. +}; // class monitor -}; // class amdsmi +#endif -/** - * \brief Namespace for flag definitions and utilities. - */ +/// Namespace for flag definitions and utilities. namespace flags { -/** - * \brief Enum representing different flags. - */ +/// Enum representing different flags. enum class Flags : uint32_t { - none = 0x0, /**< \brief No flags set */ - sync = 0x1, /**< \brief Synchronization flag */ + none = 0x0, ///< No flags set. + sync = 0x1, ///< Synchronization flag. }; -/** - * \brief Wrapper for Flags with utility operations. - */ +/// Wrapper for Flags with utility operations. struct FlagTag { - Flags value{Flags::none}; /**< \brief Underlying flag value */ + Flags value{Flags::none}; ///< Underlying flag value. - /** \brief Construct from a specific flag */ + /// Construct from a specific flag. constexpr FlagTag(Flags v) : value(v) {} - /** \brief Bitwise OR operator for combining flags */ + /// Bitwise OR operator for combining flags. friend constexpr FlagTag operator|(FlagTag a, FlagTag b) { return FlagTag( static_cast(static_cast(a.value) | static_cast(b.value))); } - /** \brief Check if a flag is set */ + /// Check if a flag is set. constexpr bool has(FlagTag f) const { return (static_cast(value) & static_cast(f.value)) != 0; @@ -939,47 +703,43 @@ struct FlagTag } // namespace flags -/** - * \brief Logger for saving benchmark results in JSON format. - * - * Handles initialization of output, storing batch data, and writing - * specialization and device information in a structured JSON file. - * - * Because the logger writes partial JSON data incrementally - * and at high volume, a JSON library is not used. - */ +/// Logger for saving benchmark results in JSON format. +/// +/// Because the logger writes partial JSON data incrementally, a JSON library is not used. class logger { public: - // Delete all copy/move constructors and assignment operators + // Delete all copy/move constructors and assignment operators. logger(const logger&) = delete; logger& operator=(const logger&) = delete; logger(logger&&) = delete; logger& operator=(logger&&) = delete; - // Singleton accessor + /// Singleton accessor. static logger& instance() { static logger instance; return instance; } - /** - * \brief Saves the program start time. - */ + /// Saves the program start time. void save_program_start_time() { m_program_start_time = std::chrono::steady_clock::now(); } - /** - * \brief Initializes the logger, and opens the output JSON and CSV files. - */ + /// Returns the program start time. + std::chrono::time_point get_program_start_time() const + { + return m_program_start_time; + } + + /// Initializes the logger, and opens the output JSON and CSV files. void init(std::string_view algorithm, size_t specialization_count, const settings& settings, flags::FlagTag flags, - const amdsmi& amdsmi) + const monitor& monitor) { m_output_batches = settings.output_batches; m_spaces_per_indent = settings.spaces_per_indent; @@ -989,7 +749,7 @@ class logger if(!m_json_out) { std::cerr << "Error: Failed to open " << settings.json_out << " for writing\n"; - std::exit(EXIT_FAILURE); + exit(EXIT_FAILURE); } if(m_outputting_csv) @@ -998,12 +758,12 @@ class logger if(!m_csv_out) { std::cerr << "Error: Failed to open " << settings.csv_out << " for writing\n"; - std::exit(EXIT_FAILURE); + exit(EXIT_FAILURE); } } m_json_out << indent( - serialize_json_prologue(algorithm, specialization_count, settings, flags, amdsmi), + serialize_json_prologue(algorithm, specialization_count, settings, flags, monitor), 0); m_json_out << "["; if(m_spaces_per_indent > 0) @@ -1021,29 +781,19 @@ class logger m_first_specialization = true; } - /** - * \brief Stores a batch of benchmark results. - * \param batch_ms Total time for the batch. - * \param iterations_ms Times for individual iterations. - * \param amdsmi_stats AMD SMI stats after the batch. - */ - void save(double batch_ms, - const std::vector& iterations_ms, - const amdsmi::stats& amdsmi_stats) + /// Stores a batch of benchmark results. + void save(double batch_ms, const std::vector& iterations_ms) { struct batch batch {}; batch.batch_ms = batch_ms; batch.iterations_ms = iterations_ms; - batch.amdsmi_stats = amdsmi_stats; m_batches.push_back(batch); } - /** - * \brief Outputs JSON and CSV specialization information. - */ + /// Outputs JSON and CSV specialization information. void output_specialization(size_t index, std::string_view name, std::string_view serialized_meta, @@ -1058,8 +808,7 @@ class logger uint16_t end_temp, double elapsed_host_secs, double elapsed_gpu_secs, - bool noise_timeout, - const amdsmi& amdsmi) + bool noise_timeout) { output_json_specialization(index, name, @@ -1075,8 +824,7 @@ class logger end_temp, elapsed_host_secs, elapsed_gpu_secs, - noise_timeout, - amdsmi); + noise_timeout); if(m_outputting_csv) { @@ -1096,9 +844,7 @@ class logger m_batches.clear(); } - /** - * \brief Outputs a summary object. - */ + /// Outputs a summary object at the end of the JSON file. void output_summary() { // Close the JSON file's outer array. @@ -1133,26 +879,24 @@ class logger struct batch { - double batch_ms; ///< Total time for the batch - std::vector iterations_ms; ///< Time per iteration - amdsmi::stats amdsmi_stats; ///< AMD SMI stats after batch + double batch_ms; ///< Total time for the batch. + std::vector iterations_ms; ///< Time per iteration. }; - /** - * \brief Serializes the start of the JSON file, - * adding the `context` object, and starting the `specializations` array. - */ + /// Serializes the start of the JSON file. + /// + /// It adds the `context` object, and adds the `specializations` key. std::string serialize_json_prologue(std::string_view algorithm, size_t specialization_count, const settings& settings, flags::FlagTag flags, - const amdsmi& amdsmi) + const monitor& monitor) { std::ostringstream ss; ss << "{"; ss << "\"context\":" - << serialize_context(algorithm, specialization_count, settings, flags, amdsmi); + << serialize_context(algorithm, specialization_count, settings, flags, monitor); ss << ","; if(m_spaces_per_indent > 0) @@ -1163,20 +907,18 @@ class logger return ss.str(); } - /** - * \brief Serializes the benchmark context into JSON. - */ + /// Serializes the benchmark context into JSON. std::string serialize_context(std::string_view algorithm, size_t specialization_count, const settings& settings, flags::FlagTag flags, - const amdsmi& amdsmi) const + const monitor& monitor) const { std::ostringstream ss; ss << "{"; - ss << "\"results_version\":\"2.0.0\""; - ss << ",\"general\":" << serialize_general(algorithm, specialization_count, amdsmi); + ss << "\"results_version\":\"3.0.0\""; + ss << ",\"general\":" << serialize_general(algorithm, specialization_count, monitor); ss << ",\"settings\":" << serialize_settings(settings); std::string custom_cli = serialize_custom_settings(settings); @@ -1186,24 +928,15 @@ class logger } ss << ",\"flags\":" << serialize_flags(flags); - if(settings.output_hip_device_properties_context) - ss << ",\"hip_device_properties\":" << serialize_hip_device_properties(); - if(settings.output_amdsmi_context) - ss << ",\"amdsmi\":" << amdsmi.serialize_context(); ss << "}"; return ss.str(); } - /** - * \brief Serializes general benchmark context info into JSON. - * - * If the macros BRANCH_NAME and/or COMMIT_HASH are defined at compile time, - * they are output as the JSON keys "branch_name" and "commit_hash" respectively. - */ + /// Serializes general benchmark context info into JSON. std::string serialize_general(std::string_view algorithm, size_t specialization_count, - const amdsmi& amdsmi) const + const monitor& monitor) const { std::ostringstream ss; ss << "{"; @@ -1211,20 +944,6 @@ class logger ss << "\"algorithm\":\"" << algorithm << "\""; ss << ",\"specialization_count\":" << specialization_count; - int device_id; - PRIMBENCH_HIP_CHECK(hipGetDevice(&device_id)); - - hipDeviceProp_t dev_prop; - PRIMBENCH_HIP_CHECK(hipGetDeviceProperties(&dev_prop, device_id)); - - ss << ",\"gpu_name\":\"" << dev_prop.name << "\""; - ss << ",\"gpu_arch\":\"" << get_arch_name(dev_prop.gcnArchName) << "\""; - - char pci_bus_id_str[32]; - PRIMBENCH_HIP_CHECK( - hipDeviceGetPCIBusId(pci_bus_id_str, sizeof(pci_bus_id_str), device_id)); - ss << ",\"gpu_pci_bus_id\":\"" << pci_bus_id_str << "\""; - #if defined(NDEBUG) const char build_type[] = "release"; #else @@ -1232,15 +951,19 @@ class logger #endif ss << ",\"library_build_type\":\"" << build_type << "\""; - ss << ",\"temp_type\":\"" << amdsmi.get_temp_type_name(amdsmi.get_temp_type()) << "\""; + ss << ",\"gpu\":" << monitor.serialize_gpu_info(); - char host_name[HOST_NAME_MAX + 1]; // +1 for null terminator + ss << ",\"backend\":" << monitor.serialize_backend_info(); + + ss << ",\"temperature_type\":\"" << monitor.get_used_temperature_type_name() << "\""; + + char host_name[HOST_NAME_MAX + 1]; // +1 for null terminator. if(gethostname(host_name, sizeof(host_name)) != 0) { std::cerr << "Error: Failed to get host name\n"; exit(EXIT_FAILURE); } - host_name[sizeof(host_name) - 1] = '\0'; // Ensure null termination + host_name[sizeof(host_name) - 1] = '\0'; // Ensure null termination. ss << ",\"host_name\":\"" << host_name << "\""; ss << ",\"date\":\"" << date() << "\""; @@ -1252,38 +975,20 @@ class logger ss << ",\"commit_hash\":\"" << COMMIT_HASH << "\""; #endif - ss << ",\"hip_version\":\"" << HIP_VERSION_MAJOR << "." << HIP_VERSION_MINOR << "." - << HIP_VERSION_PATCH << "-" << HIP_VERSION_GITHASH << "\""; - ss << ",\"clang_version\":\"" << __clang_version__ << "\""; - ss << "}"; return ss.str(); } - /** - * \brief Extracts the GPU architecture name (gfx123) from a full gcnArchName string. - */ - std::string_view get_arch_name(std::string_view gcn_arch_name) const - { - // Find the position of the first ':' (if any) - size_t pos = gcn_arch_name.find(':'); - if(pos == std::string_view::npos) - return gcn_arch_name; // no colon, return the whole string - return gcn_arch_name.substr(0, pos); // return only the part before ':' - } - - /** - * \brief Returns the local date and time as an RFC3339 string (yyyy-mm-ddTHH:MM:SS±HH:MM). - */ + /// Returns the local date and time as an RFC3339 string: `yyyy-mm-ddTHH:MM:SS±HH:MM`. std::string date() const { using namespace std::chrono; - // Get current time as system_clock::time_point + // Get current time as system_clock::time_point. auto now = system_clock::now(); std::time_t now_c = system_clock::to_time_t(now); - // Convert to local time + // Convert to local time. std::tm local_tm{}; #if defined(_WIN32) localtime_s(&local_tm, &now_c); @@ -1291,11 +996,11 @@ class logger localtime_r(&now_c, &local_tm); #endif - // Format date and time + // Format date and time. std::ostringstream oss; oss << std::put_time(&local_tm, "%Y-%m-%dT%H:%M:%S"); - // Compute timezone offset + // Compute timezone offset. std::tm utc_tm{}; #if defined(_WIN32) gmtime_s(&utc_tm, &now_c); @@ -1303,7 +1008,7 @@ class logger gmtime_r(&now_c, &utc_tm); #endif - // Offset in seconds + // Offset in seconds. int offset_sec = static_cast(std::difftime(std::mktime(&local_tm), std::mktime(&utc_tm))); char sign = offset_sec >= 0 ? '+' : '-'; @@ -1311,16 +1016,14 @@ class logger int offset_hours = offset_sec / 3600; int offset_minutes = (offset_sec % 3600) / 60; - // Append timezone + // Append timezone. oss << sign << std::setw(2) << std::setfill('0') << offset_hours << ':' << std::setw(2) << std::setfill('0') << offset_minutes; return oss.str(); } - /** - * \brief Serializes CLI settings into JSON. - */ + /// Serializes CLI settings into JSON. std::string serialize_settings(const settings& settings) const { std::ostringstream ss; @@ -1333,7 +1036,7 @@ class logger ss << ",\"seed\":" << s.seed; ss << ",\"json_out\":\"" << s.json_out << "\""; ss << ",\"csv_out\":\"" << s.csv_out << "\""; - ss << ",\"filter\":\"" << s.filter << "\""; + ss << ",\"filter\":\"" << json_escape(s.filter) << "\""; ss << ",\"dry\":" << s.dry; ss << ",\"min_gpu_ms_per_batch\":" << s.min_gpu_ms_per_batch; ss << ",\"min_secs\":" << s.min_secs; @@ -1344,9 +1047,6 @@ class logger ss << ",\"max_gpu_temp\":" << s.max_gpu_temp; ss << ",\"max_warming_secs\":" << s.max_warming_secs; ss << ",\"max_cooling_secs\":" << s.max_cooling_secs; - ss << ",\"output_hip_device_properties_context\":" - << s.output_hip_device_properties_context; - ss << ",\"output_amdsmi_context\":" << s.output_amdsmi_context; ss << ",\"output_batches\":" << s.output_batches; ss << ",\"spaces_per_indent\":" << s.spaces_per_indent; ss << ",\"stream_blocking_timeout_secs\":" << s.stream_blocking_timeout_secs; @@ -1355,9 +1055,36 @@ class logger return ss.str(); } - /** - * \brief Serializes custom CLI settings into JSON. - */ + /// Escapes a JSON string, so it transforms "\" to "\\". + std::string json_escape(std::string_view s) const + { + std::ostringstream o; + for(unsigned char c : s) + { + switch(c) + { + case '\"': o << "\\\""; break; + case '\\': o << "\\\\"; break; + case '\b': o << "\\b"; break; + case '\f': o << "\\f"; break; + case '\n': o << "\\n"; break; + case '\r': o << "\\r"; break; + case '\t': o << "\\t"; break; + default: + if(c < 0x20) + { + o << "\\u" << std::hex << std::setw(4) << std::setfill('0') << (int)c; + } + else + { + o << c; + } + } + } + return o.str(); + } + + /// Serializes custom CLI settings into JSON. std::string serialize_custom_settings(const settings& settings) const { if(settings.custom_args.empty()) @@ -1398,9 +1125,7 @@ class logger return ss.str(); } - /** - * \brief Formats a JSON string with indentation. - */ + /// Formats a JSON string with indentation. std::string indent(std::string_view str, size_t starting_indent_level) { if(m_spaces_per_indent == 0) @@ -1419,7 +1144,7 @@ class logger if(c == '\"') { - // Detect escaped quotes + // Detect escaped quotes. bool escaped = false; size_t j = i; while(j > 0 && str[--j] == '\\') @@ -1464,9 +1189,7 @@ class logger return out.str(); } - /** - * \brief Serializes benchmark flags into JSON. - */ + /// Serializes benchmark flags into JSON. std::string serialize_flags(flags::FlagTag flags) const { std::ostringstream ss; @@ -1476,228 +1199,7 @@ class logger return ss.str(); } - /** - * \brief Returns a JSON string describing the current HIP device's properties. - */ - std::string serialize_hip_device_properties() const - { - std::ostringstream ss; - ss << "{"; - ss << std::boolalpha; - - bool first = true; - auto add_comma = [&]() - { - if(!first) - ss << ","; - first = false; - }; - - auto add_field = [&](std::string_view name, const auto& value) - { - add_comma(); - ss << "\"" << name << "\":" << value; - }; - - auto add_bool = [&](std::string_view name, const bool& value) { add_field(name, value); }; - - auto add_string = [&](std::string_view name, const std::string& string) - { - add_comma(); - ss << "\"" << name << "\":\"" << string << "\""; - }; - - auto add_dim2 = [&](std::string_view name, auto&& arr) - { - add_comma(); - ss << "\"" << name << "\":{\"x\":" << arr[0] << ",\"y\":" << arr[1] << "}"; - }; - - auto add_dim3 = [&](std::string_view name, auto&& arr) - { - add_comma(); - ss << "\"" << name << "\":{\"x\":" << arr[0] << ",\"y\":" << arr[1] - << ",\"z\":" << arr[2] << "}"; - }; - - int device_id; - PRIMBENCH_HIP_CHECK(hipGetDevice(&device_id)); - - hipDeviceProp_t dev_prop; - PRIMBENCH_HIP_CHECK(hipGetDeviceProperties(&dev_prop, device_id)); - - ss << "\"identity\":{"; - add_string("name", dev_prop.name); - add_string("gcn_arch_name", dev_prop.gcnArchName); - add_field("asic_revision", dev_prop.asicRevision); - ss << "}"; - - ss << ",\"clocks\":{"; - first = true; - add_field("core_hz", dev_prop.clockRate); - add_field("instruction_hz", dev_prop.clockInstructionRate); - add_field("memory_hz", dev_prop.memoryClockRate); - ss << "}"; - - ss << ",\"compute\":{"; - - first = true; - add_field("mode", dev_prop.computeMode); - - ss << ",\"capability\":{"; - first = true; - add_field("major", dev_prop.major); - add_field("minor", dev_prop.minor); - ss << "}"; - - ss << ",\"execution\":{"; - first = true; - add_field("max_threads_per_block", dev_prop.maxThreadsPerBlock); - add_field("max_threads_per_multiprocessor", dev_prop.maxThreadsPerMultiProcessor); - add_field("multi_processor_count", dev_prop.multiProcessorCount); - add_field("regs_per_block", dev_prop.regsPerBlock); - add_field("warp_size", dev_prop.warpSize); - ss << "}"; - - ss << ",\"limits\":{"; - first = true; - add_dim3("grid_size", dev_prop.maxGridSize); - add_dim3("threads_dim", dev_prop.maxThreadsDim); - ss << "}"; - - ss << "}"; // End of "compute" object. - - ss << ",\"memory\":{"; - - ss << "\"global\":{"; - first = true; - add_field("total", dev_prop.totalGlobalMem); - add_field("pitch", dev_prop.memPitch); - add_field("bus_width", dev_prop.memoryBusWidth); - add_field("l2_cache_size", dev_prop.l2CacheSize); - ss << "}"; - - ss << ",\"shared\":{"; - first = true; - add_field("per_block", dev_prop.sharedMemPerBlock); - add_field("per_multi_processor", dev_prop.maxSharedMemoryPerMultiProcessor); - ss << "}"; - - ss << ",\"const\":{"; - first = true; - add_field("total", dev_prop.totalConstMem); - ss << "}"; - - ss << "}"; // End of "memory" object. - - ss << ",\"pci\":{"; - first = true; - add_field("bus_id", dev_prop.pciBusID); - add_field("device_id", dev_prop.pciDeviceID); - add_field("domain_id", dev_prop.pciDomainID); - ss << "}"; - - ss << ",\"texture\":{"; - first = true; - add_field("alignment", dev_prop.textureAlignment); - add_field("pitch_alignment", dev_prop.texturePitchAlignment); - add_field("max_1d", dev_prop.maxTexture1D); - add_field("max_1d_linear", dev_prop.maxTexture1DLinear); - add_dim2("max_2d", dev_prop.maxTexture2D); - add_dim3("max_3d", dev_prop.maxTexture3D); - ss << "}"; - - ss << ",\"capabilities\":{"; - - ss << "\"architecture\":{"; - - const auto arch = dev_prop.arch; - - ss << "\"atomics\":{"; - first = true; - add_bool("float_atomic_add", arch.hasFloatAtomicAdd); - add_bool("global_float_atomic_exch", arch.hasGlobalFloatAtomicExch); - add_bool("global_int32_atomics", arch.hasGlobalInt32Atomics); - add_bool("global_int64_atomics", arch.hasGlobalInt64Atomics); - add_bool("shared_float_atomic_exch", arch.hasSharedFloatAtomicExch); - add_bool("shared_int32_atomics", arch.hasSharedInt32Atomics); - add_bool("shared_int64_atomics", arch.hasSharedInt64Atomics); - ss << "}"; - - ss << ",\"warp\":{"; - first = true; - add_bool("warp_ballot", arch.hasWarpBallot); - add_bool("warp_shuffle", arch.hasWarpShuffle); - add_bool("warp_vote", arch.hasWarpVote); - ss << "}"; - - ss << ",\"misc\":{"; - first = true; - add_bool("3d_grid", arch.has3dGrid); - add_bool("doubles", arch.hasDoubles); - add_bool("dynamic_parallelism", arch.hasDynamicParallelism); - add_bool("funnel_shift", arch.hasFunnelShift); - add_bool("surface_funcs", arch.hasSurfaceFuncs); - add_bool("sync_threads_ext", arch.hasSyncThreadsExt); - add_bool("thread_fence_system", arch.hasThreadFenceSystem); - ss << "}"; - - ss << "}"; // End of "architecture" object. - - ss << ",\"cooperative_multi_device\":{"; - - ss << "\"launch\":" << static_cast(dev_prop.cooperativeMultiDeviceLaunch); - - ss << ",\"unmatched\":{"; - first = true; - add_bool("block_dim", dev_prop.cooperativeMultiDeviceUnmatchedBlockDim); - add_bool("func", dev_prop.cooperativeMultiDeviceUnmatchedFunc); - add_bool("grid_dim", dev_prop.cooperativeMultiDeviceUnmatchedGridDim); - add_bool("shared_mem", dev_prop.cooperativeMultiDeviceUnmatchedSharedMem); - ss << "}"; - - ss << "}"; // End of "cooperative_multi_device" object. - - ss << ",\"device\":{"; - - ss << "\"compute\":{"; - first = true; - add_bool("concurrent_kernels", dev_prop.concurrentKernels); - add_bool("cooperative_launch", dev_prop.cooperativeLaunch); - add_bool("is_large_bar", dev_prop.isLargeBar); - add_bool("kernel_exec_timeout_enabled", dev_prop.kernelExecTimeoutEnabled); - ss << "}"; - - ss << ",\"memory\":{"; - first = true; - add_bool("can_map_host_memory", dev_prop.canMapHostMemory); - add_bool("concurrent_managed_access", dev_prop.concurrentManagedAccess); - add_bool("direct_managed_mem_access_from_host", dev_prop.directManagedMemAccessFromHost); - add_bool("managed_memory", dev_prop.managedMemory); - add_bool("pageable_memory_access", dev_prop.pageableMemoryAccess); - add_bool("pageable_memory_access_uses_host_page_tables", - dev_prop.pageableMemoryAccessUsesHostPageTables); - ss << "}"; - - ss << ",\"misc\":{"; - first = true; - add_bool("ecc_enabled", dev_prop.ECCEnabled); - add_bool("integrated", dev_prop.integrated); - add_bool("is_multi_gpu_board", dev_prop.isMultiGpuBoard); - add_bool("tcc_driver", dev_prop.tccDriver); - ss << "}"; - - ss << "}"; // End of "device" object. - - ss << "}"; // End of "capabilities" object. - - ss << "}"; - return ss.str(); - } - - /** - * \brief Outputs specialization information to JSON. - */ + /// Outputs specialization information to JSON. void output_json_specialization(size_t index, std::string_view name, std::string_view serialized_meta, @@ -1712,8 +1214,7 @@ class logger uint16_t end_temp, double elapsed_host_secs, double elapsed_gpu_secs, - bool noise_timeout, - const amdsmi& amdsmi) + bool noise_timeout) { // Specializations need a comma between them. if(!m_first_specialization) @@ -1739,16 +1240,13 @@ class logger end_temp, elapsed_host_secs, elapsed_gpu_secs, - noise_timeout, - amdsmi), + noise_timeout), 2); m_json_out.flush(); } - /** - * \brief Outputs specialization information to CSV. - */ + /// Outputs specialization information to CSV. void output_csv_specialization(size_t index, std::string_view name, double bytes_per_sec, @@ -1762,9 +1260,7 @@ class logger << std::flush; } - /** - * \brief Serializes a specialization into JSON format. - */ + /// Serializes a specialization into JSON format. std::string serialize_specialization(size_t index, std::string_view name, std::string_view serialized_meta, @@ -1779,8 +1275,7 @@ class logger uint16_t end_temp, double elapsed_host_secs, double elapsed_gpu_secs, - bool noise_timeout, - const amdsmi& amdsmi) const + bool noise_timeout) const { std::ostringstream ss; ss << "{"; @@ -1827,8 +1322,6 @@ class logger ss << "{"; ss << "\"batch_ms\":" << batch.batch_ms; ss << ",\"iterations_ms\":" << serialize_iterations_ms(batch.iterations_ms); - ss << ",\"amdsmi_stats_after_iterations\":" - << amdsmi.serialize_stats(batch.amdsmi_stats); ss << "}"; } ss << "]"; @@ -1838,9 +1331,7 @@ class logger return ss.str(); } - /** - * \brief Serializes iteration times into JSON array. - */ + /// Serializes iteration times into JSON array. std::string serialize_iterations_ms(const std::vector& iterations_ms) const { std::ostringstream ss; @@ -1855,9 +1346,7 @@ class logger return ss.str(); } - /** - * \brief Serializes summary info into JSON format. - */ + /// Serializes summary info into JSON format. std::string serialize_summary() const { std::ostringstream ss; @@ -1881,146 +1370,166 @@ class logger std::chrono::time_point m_program_start_time; - std::ofstream m_json_out; ///< JSON output file stream - std::ofstream m_csv_out; ///< CSV output file stream - bool m_first_specialization; ///< True if first specialization output - std::vector m_batches; ///< Stored batch results - bool m_output_batches; ///< Whether to output each batch - uint32_t m_spaces_per_indent; ///< JSON indentation spaces - double m_total_elapsed_gpu_secs = 0; ///< Number of elapsed GPU seconds - uint32_t m_noise_timeouts = 0; ///< Number of noise timeouts - bool m_outputting_csv; ///< Whether a CSV file is output + std::ofstream m_json_out; ///< JSON output file stream. + std::ofstream m_csv_out; ///< CSV output file stream. + bool m_first_specialization; ///< True if first specialization output. + std::vector m_batches; ///< Stored batch results. + bool m_output_batches; ///< Whether to output each batch. + uint32_t m_spaces_per_indent; ///< JSON indentation spaces. + double m_total_elapsed_gpu_secs = 0; ///< Number of elapsed GPU seconds. + uint32_t m_noise_timeouts = 0; ///< Number of noise timeouts. + bool m_outputting_csv; ///< Whether a CSV file is output. }; // class logger -/** - * \brief Provides functions for progress display and formatting during GPU benchmarking. - */ +/// Provides functions for progress display and formatting during GPU benchmarking. namespace progress { -/// \brief Column widths and formatting constants. -static constexpr int noise_col_width = 5; -static constexpr int gpu_temp_col_width = 6; -static constexpr int bytes_per_sec_col_width = 9; -static constexpr const char* horizontal_bar = u8"─"; - -/** - * \brief Prints the table header for dry-run mode output. - * - * \param algo_name Algorithm name to display. - * \param spec_col_width Width of the specialization column. - * \param family_col_width Width of the family column. - * \param specialization_count Total number of specializations. - */ +// Column widths and formatting constants. +static constexpr int noise_column_width = 5; +static constexpr int gpu_temp_column_width = 6; +static constexpr int bytes_per_sec_column_width = 9; +static constexpr const char* horizontal_bar = u8"─"; + +/// Prints the table header for dry-run mode output. inline void print_dry_header(std::string_view algo_name, - size_t spec_col_width, - size_t family_col_width, + size_t specialization_column_width, + size_t index_column_width, size_t specialization_count) { - std::string status_header = "Status of " + std::string(algo_name); - size_t status_col_width = status_header.size(); + std::string status_header = "Status of " + std::string(algo_name); + size_t status_column_width = status_header.size(); - std::cout << std::setw(status_col_width) << std::left << status_header << " " - << std::setw(spec_col_width) << std::left << "Specialization" << " " << "Index/" - << specialization_count << "\n"; + std::cout << std::setw(status_column_width) << std::left << status_header << " " + << std::setw(specialization_column_width) << std::left << "Specialization"; - size_t underline_width = status_col_width + 2 + spec_col_width + 2 + family_col_width; + if(specialization_count > 1) + { + std::cout << " " << std::setw(index_column_width) << std::right + << ("Index/" + std::to_string(specialization_count)); + } + std::cout << "\n"; + + size_t underline_width = status_column_width + 2 + specialization_column_width; + + if(specialization_count > 1) + { + underline_width += 2 + index_column_width; + } for(size_t i = 0; i < underline_width; ++i) std::cout << horizontal_bar; std::cout << "\n"; } -/** - * \brief Prints the table header for algorithm progress output. - * - * \param algo_name Algorithm name to display. - * \param spec_col_width Width of the specialization column. - * \param family_col_width Width of the family column. - * \param specialization_count Total number of specializations. - * \param noise_timeout_secs Duration before noisy timeout. - */ +/// Prints the table header for algorithm progress output. inline void print_header(std::string_view algo_name, - size_t spec_col_width, - size_t family_col_width, + size_t specialization_column_width, + size_t index_column_width, size_t specialization_count, std::chrono::seconds::rep noise_timeout_secs) { std::string status_header = "Status of " + std::string(algo_name); std::string noisy_status = "Noisy timed out after " + std::to_string(noise_timeout_secs) + "s"; - size_t status_col_width = std::max(status_header.size(), noisy_status.size()); + size_t status_column_width = std::max(status_header.size(), noisy_status.size()); - std::cout << std::setw(status_col_width) << std::left << status_header << " " - << std::setw(noise_col_width) << std::left << "Noise" << " " - << std::setw(gpu_temp_col_width) << std::left << "GPU °C" << " " - << std::setw(bytes_per_sec_col_width) << std::left << "Bytes/sec" << " " - << std::setw(spec_col_width) << std::left << "Specialization" << " " << "Index/" - << specialization_count << "\n"; + std::cout << std::setw(status_column_width) << std::left << status_header << " " + << std::setw(noise_column_width) << std::left << "Noise" << " " + << std::setw(gpu_temp_column_width) << std::left << "GPU °C" << " " + << std::setw(bytes_per_sec_column_width) << std::left << "Bytes/sec" << " " + << std::setw(specialization_column_width) << std::left << "Specialization"; - size_t underline_width = status_col_width + 2 + noise_col_width + 2 + gpu_temp_col_width + 2 - + bytes_per_sec_col_width + 2 + spec_col_width + 2 + family_col_width; + if(specialization_count > 1) + { + std::cout << " " << std::setw(index_column_width) << std::right + << ("Index/" + std::to_string(specialization_count)); + } + std::cout << "\n"; + + size_t underline_width = status_column_width + 2 + noise_column_width + 2 + + gpu_temp_column_width + 2 + bytes_per_sec_column_width + 2 + + specialization_column_width; + + if(specialization_count > 1) + { + underline_width += 2 + index_column_width; + } for(size_t i = 0; i < underline_width; ++i) std::cout << horizontal_bar; std::cout << "\n"; } -/** - * \brief Displays GPU warming progress. - * - * \param gpu_temp Current GPU temperature. - * \param min_gpu_temp Minimum temperature target. - */ +/// Displays GPU warming progress. inline void print_warming(uint16_t gpu_temp, uint16_t min_gpu_temp) { std::cout << clearline << "Warming GPU from " << blue << gpu_temp << "°C" << reset << " to " << green << min_gpu_temp << "°C" << reset << std::flush; } -/** - * \brief Displays GPU cooling progress. - * - * \param gpu_temp Current GPU temperature. - * \param max_gpu_temp Maximum temperature target. - */ +/// Displays GPU cooling progress. inline void print_cooling(uint16_t gpu_temp, uint16_t max_gpu_temp) { std::cout << clearline << "Cooling GPU from " << red << gpu_temp << "°C" << reset << " to " << green << max_gpu_temp << "°C" << reset << std::flush; } -/** - * \brief Prints a single line of output for dry-run mode. - * - * Displays only status, specialization and family index. - * Used when simulating algorithm execution without actually running benchmarks. - */ +/// Prints a single line of output for dry-run mode, containing the status, specialization and index. inline void print_dry_progress(std::string_view specialization, std::string_view algo_name, - size_t family_index, - size_t spec_col_width, - size_t family_col_width) + size_t specialization_index, + size_t specialization_column_width, + size_t index_column_width, + bool print_index) { std::ostringstream line; - std::string status_header = "Status of " + std::string(algo_name); - size_t status_col_width = status_header.size(); + std::string status_header = "Status of " + std::string(algo_name); + size_t status_column_width = status_header.size(); - line << clearline << green << std::setw(status_col_width) << std::left << "Success" << reset; + line << clearline << green << std::setw(status_column_width) << std::left << "Success" << reset; - line << " " << std::setw(spec_col_width) << std::left << specialization; - line << " " << std::setw(family_col_width) << std::right << family_index; + line << " " << std::setw(specialization_column_width) << std::left << specialization; + + if(print_index) + { + line << " " << std::setw(index_column_width) << std::right << specialization_index; + } std::cout << line.str() << "\n" << std::flush; } -/** - * \brief Prints real-time progress updates for algorithm execution. - * - * Displays bytes per second, temperature, and specialization data. - * Highlights noisy or timed-out iterations with color-coded output. - */ +/// Returns an estimated time in the format "9h 59m 59s". +inline std::string format_eta(double remaining_secs) +{ + assert(remaining_secs > 0.0); + + if(remaining_secs > 35999.0) + remaining_secs = 35999.0; // 9h 59m 59s + + int h = static_cast(remaining_secs) / 3600; + int m = (static_cast(remaining_secs) % 3600) / 60; + int s = static_cast(remaining_secs) % 60; + + std::ostringstream oss; + if(h > 0) + { + oss << h << "h " << m << "m " << s << "s"; + } + else if(m > 0) + { + oss << std::setw(3) << "" << m << "m " << s << "s"; + } + else + { + oss << std::setw(7) << "" << s << "s"; + } + + return oss.str(); +} + +/// Prints real-time progress updates for algorithm execution. inline void print_progress(uint64_t iteration, double noise_percent, double bytes_per_sec, @@ -2028,35 +1537,37 @@ inline void print_progress(uint64_t iteration, std::string_view specialization, std::string_view algo_name, uint64_t batch_window_size, - size_t family_index, - size_t spec_col_width, - size_t family_col_width, + size_t specialization_index, + size_t specialization_column_width, + size_t index_column_width, + bool print_index, double elapsed_host_secs, double noise_timeout_secs, double noise_tolerance_percent, - uint16_t gpu_temp) + uint16_t gpu_temp, + double estimated_remaining_secs) { std::string status_header = "Status of " + std::string(algo_name); std::chrono::seconds::rep secs = noise_timeout_secs; // Casts to an integer. std::string noisy_status = "Noisy timed out after " + std::to_string(secs) + "s"; - size_t status_col_width = std::max(status_header.size(), noisy_status.size()); + size_t status_column_width = std::max(status_header.size(), noisy_status.size()); std::string batch_str = status_msg.empty() ? "Batch " + std::to_string(iteration) + "/" + std::to_string(batch_window_size) : std::string(status_msg); size_t bar_width = 0; - if(batch_str.size() + 1 < status_col_width) + if(batch_str.size() + 1 < status_column_width) { - bar_width = status_col_width - batch_str.size() - 1; + bar_width = status_column_width - batch_str.size() - 1; } std::ostringstream line; line << clearline << batch_str; - // Progress bar (only shown during iteration) + // Progress bar (only shown during iteration). if(status_msg.empty()) { line << " "; @@ -2064,7 +1575,7 @@ inline void print_progress(uint64_t iteration, uint64_t capped_iteration = std::min(iteration, batch_window_size); size_t filled = (bar_width * capped_iteration) / batch_window_size; - // Compute fraction of the yellow noise timeout overlay + // Compute fraction of the yellow noise timeout overlay. size_t yellow_chars = 0; if(filled >= bar_width) { @@ -2072,7 +1583,7 @@ inline void print_progress(uint64_t iteration, yellow_chars = bar_width * frac; } - // Build the progress bar in one pass + // Build the progress bar in one pass. for(size_t j = 0; j < bar_width; ++j) { if(j < yellow_chars) @@ -2085,12 +1596,12 @@ inline void print_progress(uint64_t iteration, line << reset; } - // Alignment for subsequent columns + // Alignment for subsequent columns. size_t used = batch_str.size() + status_msg.empty() + (status_msg.empty() ? bar_width : 0); - if(used < status_col_width) - line << std::string(status_col_width - used, ' '); + if(used < status_column_width) + line << std::string(status_column_width - used, ' '); - // Noise % + // Noise %. std::ostringstream percent_stream; if(noise_percent >= 100.0) @@ -2104,21 +1615,34 @@ inline void print_progress(uint64_t iteration, auto& noise_color = (noise_percent > noise_tolerance_percent) ? red : reset; - line << " " << noise_color << std::setw(noise_col_width - sizeof('%')) << std::right + line << " " << noise_color << std::setw(noise_column_width - sizeof('%')) << std::right << percent_stream.str() << "%" << reset; - // GPU temperature - line << " " << std::setw(gpu_temp_col_width) << std::right << gpu_temp; + // GPU temperature. + line << " " << std::setw(gpu_temp_column_width) << std::right << gpu_temp; - // Bytes/sec - line << " " << std::setw(bytes_per_sec_col_width) << std::right << std::scientific + // Bytes/sec. + line << " " << std::setw(bytes_per_sec_column_width) << std::right << std::scientific << std::setprecision(2) << bytes_per_sec; - // Specialization and index - line << " " << std::setw(spec_col_width) << std::left << specialization; - line << " " << std::setw(family_col_width) << std::right << family_index; + // Specialization. + line << " " << std::setw(specialization_column_width) << std::left << specialization; - // Colorized status messages + // Index. + if(print_index) + { + if(status_msg.empty() && estimated_remaining_secs > 0.0) + { + line << " " << std::setw(index_column_width) << std::right + << format_eta(estimated_remaining_secs); + } + else + { + line << " " << std::setw(index_column_width) << std::right << specialization_index; + } + } + + // Colorized status messages. if(status_msg.find("Success") != std::string::npos) std::cout << green; else if(status_msg.find("Noisy timed out") != std::string::npos) @@ -2126,68 +1650,107 @@ inline void print_progress(uint64_t iteration, std::cout << line.str() << std::flush; } + } // namespace progress -/** - * \brief Manages synchronization between host and GPU streams by blocking execution - * until explicitly unblocked or a timeout occurs. - */ +#ifdef __HIP__ + +/// Kernel that blocks the GPU stream until unblocked or timeout occurs. +__global__ +void block_stream_kernel(volatile int32_t* is_blocked, + volatile int32_t* timeout_flag, + double timeout_seconds, + long long int wall_clock_rate) +{ + const long long int start_time = wall_clock64(); + const long long int timeout_cycles + = static_cast(timeout_seconds * wall_clock_rate * 1000.0); + + while(*is_blocked == 1) + { + if(wall_clock64() - start_time > timeout_cycles) + { + *timeout_flag = 1; // Signal timeout to host. + break; // Exit loop. + } + } +} + +#elif defined(__CUDACC__) + +/// Kernel that blocks the GPU stream until unblocked or timeout occurs. +__global__ +void block_stream_kernel(volatile int32_t* is_blocked, + volatile int32_t* timeout_flag, + double timeout_seconds) +{ + using clock = cuda::std::chrono::high_resolution_clock; + auto start_time = clock::now(); + cuda::std::chrono::duration timeout_duration(timeout_seconds); + + while(*is_blocked == 1) + { + auto elapsed = cuda::std::chrono::duration(clock::now() - start_time); + if(elapsed >= timeout_duration) + { + *timeout_flag = 1; // Signal timeout to host. + break; // Exit loop. + } + } +} + +#endif + +/// Manages synchronization between host and GPU streams +/// by blocking execution, until explicitly unblocked or a timeout occurs. class stream_blocker { public: stream_blocker() = delete; - /** - * \brief Constructs a stream_blocker for a given HIP stream. - */ - stream_blocker(hipStream_t stream, double stream_blocking_timeout_secs) + /// Constructs a stream_blocker for a given stream. + stream_blocker(stream_t stream, double stream_blocking_timeout_secs) : m_stream(stream), m_stream_blocking_timeout_secs(stream_blocking_timeout_secs) { - // Register host memory for blocking flags - PRIMBENCH_HIP_CHECK( - hipHostRegister(&m_host_flag, sizeof(m_host_flag), hipHostRegisterMapped)); - PRIMBENCH_HIP_CHECK(hipHostRegister(&m_host_timeout_flag, - sizeof(m_host_timeout_flag), - hipHostRegisterMapped)); - - // Get device pointers to mapped host memory using temporary non-volatile pointers + // Register host memory for blocking flags. + PRIMBENCH_CHECK(host_register(&m_host_flag, sizeof(m_host_flag), host_register_mapped)); + PRIMBENCH_CHECK( + host_register(&m_host_timeout_flag, sizeof(m_host_timeout_flag), host_register_mapped)); + + // Temporary non-volatile pointers. int32_t* temp_device_flag = nullptr; int32_t* temp_device_timeout_flag = nullptr; - PRIMBENCH_HIP_CHECK( - hipHostGetDevicePointer(reinterpret_cast(&temp_device_flag), &m_host_flag, 0)); - PRIMBENCH_HIP_CHECK( - hipHostGetDevicePointer(reinterpret_cast(&temp_device_timeout_flag), - &m_host_timeout_flag, - 0)); + // Get device pointers to mapped host memory. + PRIMBENCH_CHECK( + host_get_device_pointer(reinterpret_cast(&temp_device_flag), &m_host_flag, 0)); + PRIMBENCH_CHECK(host_get_device_pointer(reinterpret_cast(&temp_device_timeout_flag), + &m_host_timeout_flag, + 0)); - // Assign to volatile members + // Assign temporary pointers to volatile members. m_device_flag = temp_device_flag; m_device_timeout_flag = temp_device_timeout_flag; +#ifdef __HIP__ + // Query wall clock rate once (constant per device). int device_id; - PRIMBENCH_HIP_CHECK(hipGetDevice(&device_id)); - - // Query wall clock rate once (constant per device) + PRIMBENCH_CHECK(hipGetDevice(&device_id)); int wall_clk_rate_k_hz = 0; - PRIMBENCH_HIP_CHECK( + PRIMBENCH_CHECK( hipDeviceGetAttribute(&wall_clk_rate_k_hz, hipDeviceAttributeWallClockRate, device_id)); - - m_wall_clock_rate = static_cast(wall_clk_rate_k_hz); + m_wall_clock_rate = wall_clk_rate_k_hz; +#endif } - /** - * \brief Destructor that unregisters host memory. - */ + /// Destructor that unregisters host memory. ~stream_blocker() { - PRIMBENCH_HIP_CHECK(hipHostUnregister(&m_host_flag)); - PRIMBENCH_HIP_CHECK(hipHostUnregister(&m_host_timeout_flag)); + PRIMBENCH_CHECK(host_unregister(&m_host_flag)); + PRIMBENCH_CHECK(host_unregister(&m_host_timeout_flag)); } - /** - * \brief Launches a blocking kernel on the stream until unblocked or timed out. - */ + /// Launches a blocking kernel on the stream until unblocked or timed out. void block() { volatile int32_t& flag = m_host_flag; @@ -2196,25 +1759,28 @@ class stream_blocker volatile int32_t& timeout_flag = m_host_timeout_flag; timeout_flag = 0; +#ifdef __HIP__ + block_stream_kernel<<>>(m_device_flag, + m_device_timeout_flag, + m_stream_blocking_timeout_secs, + m_wall_clock_rate); +#elif defined(__CUDACC__) block_stream_kernel<<>>(m_device_flag, m_device_timeout_flag, - m_wall_clock_rate, m_stream_blocking_timeout_secs); +#endif } - /** - * \brief Unblocks the stream by resetting the blocking flag. - */ + /// Unblocks the stream by resetting the blocking flag. void unblock() { volatile int32_t& flag = m_host_flag; flag = 0; } - /** - * \brief Checks if the GPU timed out during blocking. - * \note Should be called after stream synchronization. - */ + /// Checks if the GPU timed out during blocking. + /// + /// This should be called after stream synchronization. void check_timeout() { if(m_host_timeout_flag) @@ -2229,7 +1795,7 @@ class stream_blocker } private: - hipStream_t m_stream; + stream_t m_stream; double m_stream_blocking_timeout_secs; long long int m_wall_clock_rate; @@ -2238,41 +1804,12 @@ class stream_blocker volatile int32_t* m_device_flag = nullptr; volatile int32_t* m_device_timeout_flag = nullptr; - - /** - * \brief Kernel that blocks the GPU stream until unblocked or timeout occurs. - * \param is_blocked Pointer to blocking flag. - * \param timeout_flag Pointer to timeout flag. - * \param wall_clock_rate Wall clock rate in kHz. - * \param timeout_seconds Timeout duration in seconds. - */ - static __global__ - void block_stream_kernel(volatile int32_t* is_blocked, - volatile int32_t* timeout_flag, - long long int wall_clock_rate, - double timeout_seconds) - { - const long long int start_time = wall_clock64(); - const long long int timeout_cycles - = static_cast(timeout_seconds * wall_clock_rate * 1000.0); - - while(*is_blocked == 1) - { - if(wall_clock64() - start_time > timeout_cycles) - { - *timeout_flag = 1; // Signal timeout to host - break; // Exit loop - } - } - } }; // class stream_blocker -/** - * \brief Simple JSON-like container. - * - * Stores key-value pairs where values can be nested JSON objects, - * strings, integers, or doubles. Provides basic serialization to JSON and to a human-readable name. - */ +/// Simple JSON-like container. +/// +/// Stores key-value pairs where values can be nested JSON objects, +/// strings, integers, or doubles. Provides basic serialization to JSON and to a human-readable name. struct json { using key_type = std::string; @@ -2283,15 +1820,14 @@ struct json // as we want a deterministic order. using map_type = std::map; - /** - * \brief Adds a key-value pair to the JSON object. - */ + /// Adds a key-value pair to the JSON object. template json& add(std::string_view key, T value) { if constexpr(std::is_same_v) { - m_map[std::string(key)] = std::make_shared(value); // Copy nested JSON. + // Copy nested JSON. + m_map[std::string(key)] = std::make_shared(value); } else { @@ -2300,9 +1836,7 @@ struct json return *this; } - /** - * \brief Serializes the JSON object to a JSON string. - */ + /// Serializes the JSON object to a JSON string. std::string serialize() const { std::ostringstream ss; @@ -2343,9 +1877,7 @@ struct json return ss.str(); } - /** - * \brief Serializes the JSON object into a human-readable name. - */ + /// Serializes the JSON object into a human-readable name. std::string serialize_name() const { static const std::vector blacklist = {"lvl", "algo"}; @@ -2399,7 +1931,8 @@ struct json } else { - // std::to_string(v) wouldn't trim trailing zeros. + // Can't use std::to_string(v), + // as that doesn't trim trailing zeros. std::ostringstream ss; ss << std::boolalpha; ss << v; @@ -2428,9 +1961,7 @@ struct json return fmt(*this, true); } - /** - * \brief Retrieves a value by key. - */ + /// Retrieves a value by key. template const T& get(std::string_view key) const { @@ -2454,17 +1985,13 @@ struct device_storage device_storage(size_t size) : m_size(size), m_device_ptr(nullptr) { if(size > 0) - { - PRIMBENCH_HIP_CHECK(hipMalloc(&m_device_ptr, size)) - } + PRIMBENCH_CHECK(gpu_malloc(&m_device_ptr, size)); } ~device_storage() { if(m_device_ptr != nullptr) - { - PRIMBENCH_HIP_CHECK(hipFree(m_device_ptr)); - } + PRIMBENCH_CHECK(gpu_free(m_device_ptr)); } size_t get_size() const @@ -2488,31 +2015,25 @@ struct device_storage void* m_device_ptr; }; // struct device_storage -/** - * \brief Cache thrashing utility to ensure the cache contains irrelevant data. - */ +/// Cache thrashing utility to ensure the cache contains irrelevant data. struct cache_thrasher { public: - /** - * \brief Initializes the cache thrasher by allocating the required memory on device. - * - * Currently, the actual largest GPU cache size cannot be queried via HIP, so this - * conservative size is used instead. Future support via HSA could make this runtime- - * queryable. - */ + /// Initializes the cache thrasher by allocating the required memory on device. + /// + /// Currently, the actual largest GPU cache size cannot be queried via HIP, so this + /// conservative size is used instead. Future support via HSA could make this runtime- + /// queryable. cache_thrasher(size_t cache_size = PRIMBENCH_GPU_CACHE_SIZE) : m_device_storage(cache_size) {} - /** - * \brief Clears the cache by thrashing memory on device. - * - * Zeros a buffer of size `m_cache_size` to evict cached data. Should be called before - * each kernel launch. - */ - void clear_cache(hipStream_t stream = hipStreamDefault) + /// Clears the cache by thrashing memory on device. + /// + /// Zeros a buffer of size `m_cache_size` to evict cached data. Should be called before + /// each kernel launch. + void clear_cache(stream_t stream) { - PRIMBENCH_HIP_CHECK( - hipMemsetAsync(m_device_storage.get_ptr(), 0, m_device_storage.get_size(), stream)); + PRIMBENCH_CHECK( + memset_async(m_device_storage.get_ptr(), 0, m_device_storage.get_size(), stream)); } // This type is not copy assignable. @@ -2524,36 +2045,38 @@ struct cache_thrasher const device_storage m_device_storage; }; // struct cache_thrasher -struct gpu_warmer +/// Warms the GPU, using complex enough dummy kernel code +/// that the compiler can't optimize it away. +__global__ +void warmup_kernel(float* data, int n) { - static constexpr int threads_per_block = 256; - static constexpr int num_items = 1 << 20; // 1 million items. - - gpu_warmer(settings& settings, amdsmi& amdsmi) - : m_settings(settings), m_amdsmi(amdsmi), m_device_storage(num_items * sizeof(float)) - {} - - static __global__ - void warmup_kernel(float* data, int n) + int idx = blockIdx.x * blockDim.x + threadIdx.x; + if(idx < n) { - int idx = blockIdx.x * blockDim.x + threadIdx.x; - if(idx < n) + float x = 0.5f + idx * 0.0001f; + + for(int i = 0; i < 10000; ++i) { - float x = 0.5f + idx * 0.0001f; + x += sinf(x) * cosf(x); + x *= 1.0000001f; + x = sqrtf(x + 1.0f); + x = logf(x + 1.0f); + } - for(int i = 0; i < 10000; ++i) - { - x += sinf(x) * cosf(x); - x *= 1.0000001f; - x = sqrtf(x + 1.0f); - x = logf(x + 1.0f); - } + data[idx] = x; + } +} + +struct gpu_warmer +{ + static constexpr int threads_per_block = 256; + static constexpr int num_items = 1 << 20; // 1 million items. - data[idx] = x; - } - } + gpu_warmer(settings& settings, monitor& monitor) + : m_settings(settings), m_monitor(monitor), m_device_storage(num_items * sizeof(float)) + {} - void warm_up(hipStream_t stream = hipStreamDefault) const + void warm_up(stream_t stream) const { auto ceil_div = [](int a, int b) -> int { return (a + b - 1) / b; }; @@ -2563,7 +2086,7 @@ struct gpu_warmer while(true) { - uint16_t gpu_temp = m_amdsmi.get_temp(); + uint16_t gpu_temp = m_monitor.get_temp(); if(gpu_temp >= s.min_gpu_temp) break; @@ -2575,7 +2098,7 @@ struct gpu_warmer warmup_kernel<<>>(m_device_storage.get_ptr(), num_items); - PRIMBENCH_HIP_CHECK(hipStreamSynchronize(stream)); + PRIMBENCH_CHECK(stream_synchronize(stream)); auto duration = std::chrono::steady_clock::now() - start; if(duration >= std::chrono::duration(s.max_warming_secs)) @@ -2596,33 +2119,27 @@ struct gpu_warmer private: const settings& m_settings; - const amdsmi& m_amdsmi; + const monitor& m_monitor; const device_storage m_device_storage; }; // struct gpu_warmer -/** - * \brief Manages benchmark execution, GPU warm-up/cool-down, timing, and logging. - * - * The `state` class coordinates benchmark runs by controlling GPU temperature, - * iteration timing, throughput calculation, and result logging. - */ +/// Manages benchmark execution, GPU warm-up/cool-down, timing, and logging. class state { public: - /** - * \brief Constructs a benchmark state. - */ state(std::string_view algo, json meta, - size_t family_index, - hipStream_t stream, + size_t specialization_index, + size_t specialization_count, + stream_t stream, logger& logger, - amdsmi& amdsmi, + monitor& monitor, stream_blocker& stream_blocker, const settings& settings, flags::FlagTag flags, - size_t spec_col_width, - size_t family_col_width, + size_t specialization_column_width, + size_t index_column_width, + bool print_index, cache_thrasher& cache, gpu_warmer& warmer) : stream(stream) @@ -2630,26 +2147,24 @@ class state , seed(settings.seed) , m_algo(algo) , m_meta(std::move(meta)) - , m_family_index(family_index) + , m_specialization_index(specialization_index) + , m_specialization_count(specialization_count) , m_logger(logger) - , m_amdsmi(amdsmi) + , m_monitor(monitor) , m_stream_blocker(stream_blocker) , m_settings(settings) , m_flags(flags) - , m_spec_col_width(spec_col_width) - , m_family_col_width(family_col_width) + , m_specialization_column_width(specialization_column_width) + , m_index_column_width(index_column_width) + , m_print_index(print_index) , m_cache(cache) , m_warmer(warmer) {} - /** - * \brief Sets the total number of items processed per iteration. - * - * This must be called exactly once before calling \ref run() or any - * memory tracking methods such as \ref add_reads() or \ref add_writes(). - * - * \param items The number of items processed per iteration. - */ + /// Sets the total number of items processed per iteration. + /// + /// This must be called exactly once before calling \ref run() or any + /// memory tracking methods such as \ref add_reads() or \ref add_writes(). void set_items(size_t items) { if(m_has_set_items) @@ -2662,19 +2177,14 @@ class state m_items = items; } - /** - * \brief Adds an estimate of global memory reads performed by the benchmark. - * - * Must be called after \ref set_items() and before any call to - * \ref add_writes(). Multiple calls accumulate total read bytes. - * - * The total number of bytes read (from all calls to this function) - * is **summed together with the total bytes written** (added via - * \ref add_writes()) to compute the reported memory throughput. - * - * \tparam T The data type of the items being read. - * \param items The number of items read. - */ + /// Adds an estimate of global memory reads performed by the benchmark. + /// + /// Must be called after \ref set_items() and before any call to + /// \ref add_writes(). Multiple calls accumulate total read bytes. + /// + /// The total number of bytes read (from all calls to this function) + /// is **summed together with the total bytes written** (added via + /// \ref add_writes()) to compute the reported memory throughput. template void add_reads(size_t items) { @@ -2693,19 +2203,14 @@ class state m_read_write_bytes += bytes; } - /** - * \brief Adds an estimate of global memory writes performed by the benchmark. - * - * Must be called after \ref set_items(). Multiple calls accumulate total - * written bytes. - * - * The total number of bytes written (from all calls to this function) - * is **summed together with the total bytes read** (added via - * \ref add_reads()) to compute the reported memory throughput. - * - * \tparam T The data type of the items being written. - * \param items The number of items written. - */ + /// Adds an estimate of global memory writes performed by the benchmark. + /// + /// Must be called after \ref set_items(). Multiple calls accumulate total + /// written bytes. + /// + /// The total number of bytes written (from all calls to this function) + /// is **summed together with the total bytes read** (added via + /// \ref add_reads()) to compute the reported memory throughput. template void add_writes(size_t items) { @@ -2720,25 +2225,20 @@ class state m_read_write_bytes += bytes; } - /** - * \brief Sets a callback to run before each iteration. - * - * Useful for resetting input data in in-place algorithms. - */ + /// Sets a callback to run before each iteration, which should be used + /// to reset the input data of in-place algorithms. void run_before_every_iteration(std::function lambda) { m_run_before_every_iteration_lambda = lambda; } - /** - * \brief Executes the benchmark loop for the provided kernel. - * - * Handles warm-up, timing, CV-based stopping, and logging. - * - * The benchmark manages all required stream synchronization internally to - * ensure accurate timing and prevent command queue buildup. Users should not - * perform any manual synchronization before or during the benchmark run. - */ + /// Executes the benchmark loop for the provided kernel. + /// + /// Handles warm-up, timing, CV-based stopping, and logging. + /// + /// The benchmark manages all required stream synchronization internally to + /// ensure accurate timing and prevent command queue buildup. Users should not + /// perform any manual synchronization before or during the benchmark run. void run(std::function kernel) { if(!m_has_set_items) @@ -2747,8 +2247,6 @@ class state exit(EXIT_FAILURE); } - const auto& s = m_settings; - std::string name = m_meta.serialize_name(); std::string serialized_meta = m_meta.serialize(); @@ -2760,138 +2258,170 @@ class state init_kernels_per_batch(kernel); // Reserve space for start and stop events for each iteration. - std::vector events(m_kernels_per_batch * 2); + std::vector events(m_kernels_per_batch * 2); for(auto& event : events) - PRIMBENCH_HIP_CHECK(hipEventCreate(&event)); + PRIMBENCH_CHECK(event_create(&event)); uint64_t iterations = 0; std::vector iterations_ms(m_kernels_per_batch); - uint16_t start_temp = m_amdsmi.get_temp(); + uint16_t start_temp = m_monitor.get_temp(); double elapsed_gpu_secs = 0.0; + double estimated_remaining_secs = get_estimated_remaining_secs(); + auto start = std::chrono::steady_clock::now(); while(true) { - iterations++; + if(run_iteration(kernel, + events, + iterations, + iterations_ms, + start, + start_temp, + elapsed_gpu_secs, + name, + serialized_meta, + bytes_per_item, + estimated_remaining_secs)) + { + break; + } + } - run_batch(events, kernel); + for(const auto& event : events) + PRIMBENCH_CHECK(event_destroy(event)); + } - fill_iterations_ms(iterations_ms, events); + // Public fields accessed directly by benchmarks. + const stream_t stream; ///< Stream used by benchmarks for kernel launches. + const size_t size; ///< Input size processed per iteration. + const uint32_t seed; ///< Random seed used for reproducible benchmark inputs. - double batch_gpu_ms = std::accumulate(iterations_ms.begin(), iterations_ms.end(), 0.0); - m_times.emplace_back(batch_gpu_ms); +private: + /// Performs a single iteration of the benchmark loop. + /// + /// @return `true` if the benchmark should stop (stable noise or timeout). + bool run_iteration(std::function& kernel, + std::vector& events, + uint64_t& iterations, + std::vector& iterations_ms, + const std::chrono::steady_clock::time_point& start, + uint16_t start_temp, + double& elapsed_gpu_secs, + const std::string& name, + const std::string& serialized_meta, + size_t bytes_per_item, + double estimated_remaining_secs) + { + const auto& s = m_settings; - amdsmi::stats amdsmi_stats = m_amdsmi.get_stats(); + iterations++; - m_logger.save(batch_gpu_ms, iterations_ms, amdsmi_stats); + run_batch(events, kernel); - // Compute noise (CV) for recent window - auto window_start = m_times.end() - std::min(iterations, s.batch_window_size); - std::vector recent_times(window_start, m_times.end()); - double recent_mean = get_mean(recent_times); - double recent_stddev = get_stddev(recent_times); - double noise_percent = get_cv(recent_times, recent_stddev, recent_mean) * 100.0; + fill_iterations_ms(iterations_ms, events); - double batch_gpu_secs = batch_gpu_ms / 1000; + double batch_gpu_ms = std::accumulate(iterations_ms.begin(), iterations_ms.end(), 0.0); + m_times.emplace_back(batch_gpu_ms); - elapsed_gpu_secs += batch_gpu_secs; + m_logger.save(batch_gpu_ms, iterations_ms); - double bytes_per_batch = m_read_write_bytes * m_kernels_per_batch; - double bytes_per_sec = bytes_per_batch / batch_gpu_secs; + // Gather batch_window_size number of iteration times. + auto window_start = m_times.end() - std::min(iterations, s.batch_window_size); + std::vector recent_times(window_start, m_times.end()); - double items_per_batch = m_items * m_kernels_per_batch; - double items_per_sec = items_per_batch / batch_gpu_secs; + // Compute noise (CV) from recent times. + double recent_mean = get_mean(recent_times); + double recent_stddev = get_stddev(recent_times); + double noise_percent = get_cv(recent_times, recent_stddev, recent_mean) * 100.0; - auto now = std::chrono::steady_clock::now(); - auto elapsed_host = now - start; + double batch_gpu_secs = batch_gpu_ms / 1000; - // Stop early if the noise stabilized. - bool stop_early = elapsed_host >= std::chrono::duration(s.min_secs) - && iterations >= s.batch_window_size - && noise_percent < s.noise_tolerance_percent; + elapsed_gpu_secs += batch_gpu_secs; - // Stop early if the benchmark has been noisy for too long. - bool noise_timeout = elapsed_host > std::chrono::duration(s.noise_timeout_secs) - && iterations >= s.batch_window_size - && noise_percent >= s.noise_tolerance_percent; + double bytes_per_batch = m_read_write_bytes * m_kernels_per_batch; + double bytes_per_sec = bytes_per_batch / batch_gpu_secs; - double elapsed_host_secs = std::chrono::duration(elapsed_host).count(); + double items_per_batch = m_items * m_kernels_per_batch; + double items_per_sec = items_per_batch / batch_gpu_secs; - std::string status; - std::chrono::seconds::rep secs = elapsed_host_secs; // Casts to an integer. - if(stop_early) - status = "Success after " + std::to_string(secs) + "s"; - else if(noise_timeout) - status = "Noisy timed out after " + std::to_string(secs) + "s"; + auto now = std::chrono::steady_clock::now(); + auto elapsed_host = now - start; - uint16_t gpu_temp = m_amdsmi.get_temp(); + // Stop early if the noise stabilized. + bool stop_early = elapsed_host >= std::chrono::duration(s.min_secs) + && iterations >= s.batch_window_size + && noise_percent < s.noise_tolerance_percent; - progress::print_progress(iterations, - noise_percent, - bytes_per_sec, - status, - name, - m_algo, - s.batch_window_size, - m_family_index, - m_spec_col_width, - m_family_col_width, - elapsed_host_secs, - s.noise_timeout_secs, - s.noise_tolerance_percent, - gpu_temp); + // Stop early if the benchmark has been noisy for too long. + bool noise_timeout = elapsed_host > std::chrono::duration(s.noise_timeout_secs) + && iterations >= s.batch_window_size + && noise_percent >= s.noise_tolerance_percent; - if(stop_early || noise_timeout) - { - std::cout << "\n"; - - m_logger.output_specialization(m_family_index, - name, - serialized_meta, - m_kernels_per_batch, - m_ms_per_batch, - bytes_per_sec, - items_per_sec, - bytes_per_item, - m_items, - noise_percent, - start_temp, - gpu_temp, - elapsed_host_secs, - elapsed_gpu_secs, - noise_timeout, - m_amdsmi); + double elapsed_host_secs = std::chrono::duration(elapsed_host).count(); - break; - } + std::string status; + std::chrono::seconds::rep secs = elapsed_host_secs; // Casts to an integer. + if(stop_early) + status = "Success after " + std::to_string(secs) + "s"; + else if(noise_timeout) + status = "Noisy timed out after " + std::to_string(secs) + "s"; + + uint16_t gpu_temp = m_monitor.get_temp(); + + progress::print_progress(iterations, + noise_percent, + bytes_per_sec, + status, + name, + m_algo, + s.batch_window_size, + m_specialization_index, + m_specialization_column_width, + m_index_column_width, + m_print_index, + elapsed_host_secs, + s.noise_timeout_secs, + s.noise_tolerance_percent, + gpu_temp, + estimated_remaining_secs); + + if(stop_early || noise_timeout) + { + std::cout << "\n"; + + m_logger.output_specialization(m_specialization_index, + name, + serialized_meta, + m_kernels_per_batch, + m_ms_per_batch, + bytes_per_sec, + items_per_sec, + bytes_per_item, + m_items, + noise_percent, + start_temp, + gpu_temp, + elapsed_host_secs, + elapsed_gpu_secs, + noise_timeout); + + return true; } - for(const auto& event : events) - PRIMBENCH_HIP_CHECK(hipEventDestroy(event)); + return false; } - /** - * \brief Public fields accessed directly by benchmarks. - */ - const hipStream_t stream; ///< HIP stream used by benchmarks for kernel launches. - const size_t size; ///< Input size processed per iteration. - const uint32_t seed; ///< Random seed used for reproducible benchmark inputs. - -private: - /** - * \brief Warms up the GPU until minimum temperature is reached. - */ + /// Warms up the GPU until minimum temperature is reached. void warm_up() const { m_warmer.warm_up(stream); } - /** - * \brief Waits for GPU to cool down below maximum temperature. - */ + /// Waits for GPU to cool down below maximum temperature. void cool_down() const { auto start = std::chrono::steady_clock::now(); @@ -2900,7 +2430,7 @@ class state while(true) { - uint16_t gpu_temp = m_amdsmi.get_temp(); + uint16_t gpu_temp = m_monitor.get_temp(); if(gpu_temp <= s.max_gpu_temp) break; @@ -2918,22 +2448,20 @@ class state } } - /** - * \brief Determines number of kernels per batch based on minimum duration. - */ + /// Determines number of kernels per batch based on minimum duration. void init_kernels_per_batch(std::function kernel) { - std::vector events(2); - std::vector iterations_ms; + std::vector events(2); + std::vector iterations_ms; m_kernels_per_batch = 1; - // Without this, the very first timed batch is very slow. + // Without this, the very first timed batch can be very slow. log("Running warmup"); for(auto& event : events) - PRIMBENCH_HIP_CHECK(hipEventCreate(&event)); + PRIMBENCH_CHECK(event_create(&event)); run_batch(events, kernel); for(const auto& event : events) - PRIMBENCH_HIP_CHECK(hipEventDestroy(event)); + PRIMBENCH_CHECK(event_destroy(event)); while(true) { @@ -2945,7 +2473,7 @@ class state iterations_ms.resize(m_kernels_per_batch); for(auto& event : events) - PRIMBENCH_HIP_CHECK(hipEventCreate(&event)); + PRIMBENCH_CHECK(event_create(&event)); run_batch(events, kernel); @@ -2955,7 +2483,7 @@ class state std::chrono::duration batch_ms(m_ms_per_batch); for(const auto& event : events) - PRIMBENCH_HIP_CHECK(hipEventDestroy(event)); + PRIMBENCH_CHECK(event_destroy(event)); if(batch_ms > std::chrono::duration(m_settings.min_gpu_ms_per_batch)) break; @@ -2965,10 +2493,8 @@ class state } } - /** - * \brief Executes a batch of kernel iterations with event timing. - */ - void run_batch(const std::vector& events, std::function kernel) + /// Executes a batch of kernel iterations with event timing. + void run_batch(const std::vector& events, std::function kernel) { for(size_t i = 0; i < m_kernels_per_batch; i++) { @@ -2978,23 +2504,24 @@ class state if(!m_settings.hot) clear_gpu_cache(stream); - // We block the stream to ensure the start event is recorded immediately before the - // kernel launch. Without this, hipEventRecord() might be queued much earlier, so the - // "start" event could capture a timestamp well before the kernel actually begins - // executing. block_stream() guarantees there is no time gap between recording the start - // event and queuing the kernel on the GPU. + // We block the stream to ensure the start event is recorded + // immediately before the kernel launch. + // Without this, event_record() might be queued much earlier, so the "start" event + // could capture a timestamp well before the kernel actually begins executing. + // block_stream() guarantees there is no time gap between recording + // the start event and queuing the kernel on the GPU. if(!m_flags.has(flags::Flags::sync)) m_stream_blocker.block(); // Even events record the start time. - PRIMBENCH_HIP_CHECK(hipEventRecord(events[i * 2], stream)); + PRIMBENCH_CHECK(event_record(events[i * 2], stream)); // In order for the event timing to be accurate, this kernel lambda // shouldn't do more than just calling the __global__ kernel function. kernel(); // Odd events record the stop time. - PRIMBENCH_HIP_CHECK(hipEventRecord(events[i * 2 + 1], stream)); + PRIMBENCH_CHECK(event_record(events[i * 2 + 1], stream)); // Allows the GPU to start running its queued events. if(!m_flags.has(flags::Flags::sync)) @@ -3004,8 +2531,8 @@ class state // We deliberately don't do this right after the kernel() call, // since that'd keep the GPU blocked for slightly longer. // The kernel lambda is still responsible for catching - // host-side algorithm errors using PRIMBENCH_HIP_CHECK(). - PRIMBENCH_HIP_CHECK(hipGetLastError()); + // host-side algorithm errors, using say PRIMBENCH_CHECK(). + PRIMBENCH_CHECK(get_last_error()); // Periodically sync to avoid overflowing the stream's command queue. // Without this, too many pending events can exhaust driver resources and cause a hang. @@ -3013,9 +2540,9 @@ class state constexpr size_t n = 64; if(i % n == n - 1) { - PRIMBENCH_HIP_CHECK(hipStreamSynchronize(stream)); + PRIMBENCH_CHECK(stream_synchronize(stream)); // Catch runtime/device execution errors. - PRIMBENCH_HIP_CHECK(hipGetLastError()); + PRIMBENCH_CHECK(get_last_error()); // Check for blocking kernel timeout after sync. if(!m_flags.has(flags::Flags::sync)) @@ -3024,52 +2551,66 @@ class state } // Final stream synchronization. - PRIMBENCH_HIP_CHECK(hipStreamSynchronize(stream)); + PRIMBENCH_CHECK(stream_synchronize(stream)); // Catch any runtime/device errors from the last kernels. - PRIMBENCH_HIP_CHECK(hipGetLastError()); + PRIMBENCH_CHECK(get_last_error()); // Final blocking kernel timeout check. if(!m_flags.has(flags::Flags::sync)) m_stream_blocker.check_timeout(); } - /** - * \brief Fills iteration times (ms) using HIP event timing. - */ - void fill_iterations_ms(std::vector& iterations_ms, - const std::vector& events) const + /// Fills iteration times (ms) using HIP event timing. + void fill_iterations_ms(std::vector& iterations_ms, + const std::vector& events) const { for(size_t i = 0; i < m_kernels_per_batch; i++) { float iteration_ms; // Gets the number of milliseconds between the start and stop event. - PRIMBENCH_HIP_CHECK( - hipEventElapsedTime(&iteration_ms, events[i * 2], events[i * 2 + 1])); + PRIMBENCH_CHECK(event_elapsed_time(&iteration_ms, events[i * 2], events[i * 2 + 1])); iterations_ms[i] = iteration_ms; } } - /** - * \brief Clears GPU caches. - */ - void clear_gpu_cache(hipStream_t stream) const + /// Clears GPU caches. + void clear_gpu_cache(stream_t stream) const { m_cache.clear_cache(stream); } - /** - * \brief Computes mean of time samples. - */ + /// Returns the estimated time remaining for the entire program. + double get_estimated_remaining_secs() const + { + // We can only estimate once at least one specialization has finished. + if(m_specialization_index == 0) + { + return 0.0; + } + + auto now = std::chrono::steady_clock::now(); + auto start = m_logger.get_program_start_time(); + + auto elapsed_total = std::chrono::duration(now - start).count(); + + // Average time taken per completed specialization. + double avg_time_per_spec = elapsed_total / m_specialization_index; + + // Number of specializations yet to be completed (including the current one). + size_t remaining_specs = m_specialization_count - m_specialization_index; + + return avg_time_per_spec * remaining_specs; + } + + /// Computes mean of time samples. double get_mean(const std::vector& times) const { return std::reduce(times.begin(), times.end()) / times.size(); } - /** - * \brief Computes standard deviation of time samples. - */ + /// Computes standard deviation of time samples. double get_stddev(const std::vector& times) const { const size_t n = times.size(); @@ -3084,9 +2625,7 @@ class state return std::sqrt(sum_sq / (n - 1)); } - /** - * \brief Computes coefficient of variation (CV). - */ + /// Computes coefficient of variation (CV). double get_cv(const std::vector& times, double stddev, double mean) const { return times.size() >= 2 ? stddev / mean : 0.0; @@ -3094,18 +2633,22 @@ class state std::string m_algo; const json m_meta; - size_t m_family_index; + + size_t m_specialization_index; + size_t m_specialization_count; logger& m_logger; - amdsmi& m_amdsmi; + monitor& m_monitor; stream_blocker& m_stream_blocker; const settings& m_settings; flags::FlagTag m_flags; - size_t m_spec_col_width; - size_t m_family_col_width; + size_t m_specialization_column_width; + + size_t m_index_column_width; + bool m_print_index; cache_thrasher& m_cache; gpu_warmer& m_warmer; @@ -3121,9 +2664,7 @@ class state size_t m_read_write_bytes = 0; }; // class state -/** - * \brief Simple command-line argument parser. - */ +/// Simple command-line argument parser. class cli { public: @@ -3141,41 +2682,50 @@ class cli std::string key = arg.substr(2); std::string value = ""; - // Only consume next argument if it's a value (doesn't start with --) - if(i + 1 < argc && std::string(argv[i + 1]).rfind("--", 0) != 0) + + // Consume all consecutive arguments that don't start with `--`. + while(i + 1 < argc && std::string(argv[i + 1]).rfind("--", 0) != 0) { - value = argv[++i]; + if(!value.empty()) + value += " "; + value += argv[++i]; } + _parsed[key] = value; } } - /// \brief Gets argument value, registering it with default and description if needed. + /// Gets argument value, registering it with default and description if needed. template T get(std::string_view name, const T& default_val, std::string_view description) { std::string key{name}; - // Register if not already registered + // Register if not already registered. if(_registered.find(key) == _registered.end()) { register_description(key, description); std::ostringstream oss; - // For bools, explicitly output "true" or "false" instead of "0" or "1" + // For bools, explicitly output "true" or "false" instead of "0" or "1". if constexpr(std::is_same_v) { - if(default_val) + oss << std::boolalpha << default_val; + } + // For vectors, join elements with spaces. + else if constexpr(is_vector::value) + { + for(size_t i = 0; i < default_val.size(); ++i) { - std::cerr << "Error: Boolean flag --" << key - << " cannot be registered with a default value of true. " - << "Flags are implicitly false.\n"; - std::exit(EXIT_FAILURE); + if(i > 0) + oss << " "; + oss << default_val[i]; } - - oss << std::boolalpha; } - oss << default_val; + else + { + oss << default_val; + } _defaults[key] = oss.str(); _registered.insert(key); @@ -3187,7 +2737,7 @@ class cli if(it != _parsed.end() && !it->second.empty()) { std::cerr << "Error: Boolean flag --" << key << " does not take a value.\n"; - std::exit(EXIT_FAILURE); + exit(EXIT_FAILURE); } } } @@ -3203,7 +2753,7 @@ class cli { std::istringstream ss(default_it->second); - // For bools, use boolalpha to parse "true"/"false" + // For bools, use boolalpha to parse "true"/"false". if constexpr(std::is_same_v) { ss >> std::boolalpha >> out; @@ -3211,7 +2761,25 @@ class cli { std::cerr << "Error: Failed to parse default for --" << key << ": invalid value \"" << default_it->second << "\"\n"; - std::exit(EXIT_FAILURE); + exit(EXIT_FAILURE); + } + } + else if constexpr(is_vector::value) + { + using Elem = typename T::value_type; + std::string token; + while(ss >> token) + { + std::istringstream es(token); + Elem e{}; + es >> e; + if(!es) + { + std::cerr << "Error: Failed to parse default vector for --" << key + << ": invalid value \"" << token << "\"\n"; + exit(EXIT_FAILURE); + } + out.push_back(e); } } else @@ -3221,7 +2789,7 @@ class cli { std::cerr << "Error: Failed to parse default for --" << key << ": invalid value \"" << default_it->second << "\"\n"; - std::exit(EXIT_FAILURE); + exit(EXIT_FAILURE); } } } @@ -3233,6 +2801,34 @@ class cli { return it->second.empty() ? true : false; } + else if constexpr(is_vector::value) + { + using Elem = typename T::value_type; + T out{}; + std::istringstream ss(it->second); + std::string token; + while(ss >> token) + { + std::istringstream es(token); + Elem e{}; + if constexpr(std::is_same_v) + { + e = token; + } + else + { + es >> e; + if(!es) + { + std::cerr << "Error: Failed to parse --" << key << ": invalid value \"" + << token << "\"\n"; + exit(EXIT_FAILURE); + } + } + out.push_back(e); + } + return out; + } else { T out{}; @@ -3251,7 +2847,7 @@ class cli { std::cerr << "Error: Failed to parse --" << key << ": invalid value \"" << it->second << "\"\n"; - std::exit(EXIT_FAILURE); + exit(EXIT_FAILURE); } // When "--size" is specified, the number can optionally be suffixed with KiB/MiB/GiB. @@ -3284,12 +2880,12 @@ class cli } } - /// \brief Returns all registered arguments with their parsed values. + /// Returns all registered arguments with their parsed values. std::map get_all_custom_options() const { std::map custom_args; - // Skip built-in arguments that are already in settings + // Skip built-in arguments that are already in settings. static const std::unordered_set builtin_args = {"help", "size", @@ -3308,8 +2904,6 @@ class cli "max-gpu-temp", "max-warming-secs", "max-cooling-secs", - "output-hip-device-properties-context", - "output-amdsmi-context", "output-batches", "spaces-per-indent", "stream-blocking-timeout-secs"}; @@ -3318,10 +2912,10 @@ class cli { if(value.empty()) { - return true; // Boolean flag + return true; // Boolean flag. } - // Check for boolean strings + // Check for boolean strings. if(value == "true") { return true; @@ -3331,7 +2925,7 @@ class cli return false; } - // Try int + // Try int. try { size_t idx = 0; @@ -3344,7 +2938,7 @@ class cli catch(...) {} - // Try double + // Try double. try { size_t idx = 0; @@ -3357,14 +2951,14 @@ class cli catch(...) {} - // Fall back to string + // Fall back to string. return value; }; - // Process all registered arguments, checking both defaults and parsed values + // Process all registered arguments, checking both defaults and parsed values. std::unordered_set processed; - // First, add all parsed custom arguments + // First, add all parsed custom arguments. for(const auto& [key, value] : _parsed) { if(builtin_args.find(key) == builtin_args.end()) @@ -3374,7 +2968,7 @@ class cli } } - // Then, add defaults for custom arguments that weren't parsed + // Then, add defaults for custom arguments that weren't parsed. for(const auto& [key, default_val] : _defaults) { if(builtin_args.find(key) == builtin_args.end() @@ -3387,7 +2981,7 @@ class cli return custom_args; } - /// \brief Prints help if requested and validates all arguments are registered. + /// Prints help if requested and validates all arguments are registered. void finalize() const { possibly_print_help(); @@ -3395,7 +2989,7 @@ class cli } private: - /// \brief Prints help message and exits if --help was requested. + /// Prints help message and exits if --help was requested. void possibly_print_help() const { if(_parsed.find("help") == _parsed.end()) @@ -3423,10 +3017,10 @@ class cli std::cout << " " << desc << "\n"; } } - std::exit(EXIT_SUCCESS); + exit(EXIT_SUCCESS); } - /// \brief Validates that all parsed arguments were registered; exits if not. + /// Validates that all parsed arguments were registered; exits if not. void validate_arguments() const { for(const auto& [key, value] : _parsed) @@ -3438,27 +3032,26 @@ class cli << "To register this argument, call .get() with a type and description:\n"; std::cerr << " executor.get(\"" << key << "\", default_value, \"Description\");\n"; - std::exit(EXIT_FAILURE); + exit(EXIT_FAILURE); } } } - std::string _appname; - std::unordered_map _parsed; + std::string _appname; ///< Stores argv[0]. + std::unordered_map + _parsed; ///< Stores pairs of passed flag_name+flag_value. - // Preserves insertion order. - std::vector> _descriptions; + std::vector> _descriptions; ///< Preserves insertion order. - // Prevents duplicate descriptions being printed. - std::unordered_set _description_keys_set; + std::unordered_set + _description_keys_set; ///< Prevents duplicate descriptions being printed. - // Store string representation of default values. - std::unordered_map _defaults; + std::unordered_map + _defaults; ///< Stores string representation of default values. - // Track which arguments were registered with set(). - std::unordered_set _registered; + std::unordered_set _registered; ///< Tracks which arguments were registered. - /// \brief Registers a description for an argument; auto-adds trailing period. + /// Registers a description for an argument; auto-adds trailing period. void register_description(const std::string& key, std::string_view description) { std::string desc{description}; @@ -3472,6 +3065,15 @@ class cli _descriptions.emplace_back(key, desc); _description_keys_set.insert(key); } + + // Helper to detect vectors. + template + struct is_vector : std::false_type + {}; + + template + struct is_vector> : std::true_type + {}; }; // class cli } // namespace detail @@ -3479,35 +3081,27 @@ class cli using json = detail::json; using state = detail::state; -/** - * \brief Used to retrieve the name of a type - * that was registered with PRIMBENCH_REGISTER_TYPE(). - */ +/// Used to retrieve the name of a type that was registered with PRIMBENCH_REGISTER_TYPE(). template std::string name() { return detail::type_name::name; } -/** - * \brief Logs a gray line of text to stdout, overwriting the previous line. - * - * This function is primarily used in benchmarks to display progress or setup messages - * (for example, "Generating matrix of size 32x64"). It accepts any number of arguments - * of varying types, concatenates them, and prints them in gray text to the console. - * - * This logging is especially helpful for diagnosing **slow setup steps** and - * **slow computers**. - * - * ### Examples - * ```cpp - * primbench::log("Loading dataset..."); - * // Output: Loading dataset... - * - * primbench::log("Generating matrix of size ", 32, "x", 64); - * // Output: Generating matrix of size 32x64 - * ``` - */ +/// This function is primarily used in benchmarks to display progress or setup messages +/// (for example, "Generating matrix of size 32x64"). It accepts any number of arguments +/// of varying types, concatenates them, and prints them as a gray line. +/// +/// This logging is especially helpful for diagnosing slow setup steps. +/// +/// Examples: +/// ```cpp +/// primbench::log("Loading dataset..."); +/// // Output: Loading dataset... +/// +/// primbench::log("Generating matrix of size ", 32, "x", 64); +/// // Output: Generating matrix of size 32x64 +/// ``` template void log(Args&&... args) { @@ -3516,63 +3110,52 @@ void log(Args&&... args) std::cout << detail::reset << std::flush; } -/** - * \brief Namespace for flag definitions and utilities. - */ +/// Namespace for flag definitions and utilities. namespace flags { -/** \brief FlagTag representing no flags */ +/// FlagTag representing no flags. inline constexpr detail::flags::FlagTag none{detail::flags::Flags::none}; -/** \brief FlagTag representing the sync flag */ +/// FlagTag representing the sync flag. inline constexpr detail::flags::FlagTag sync{detail::flags::Flags::sync}; } // namespace flags -/** - * \brief Base interface for all benchmark specializations. - * - * A benchmark implementation describes: - * - the algorithm (`meta()["algo"]`), - * - a JSON-formatted specialization identifier (other keys in `meta()`), - * - and the code that performs the timed measurement (`run()`). - * - * The executor uses this interface to: - * - validate and sort benchmarks, - * - construct per-benchmark state objects, - * - run kernels and collect performance data, - * - and emit structured JSON results. - */ +/// Interface for all benchmark specializations. +/// +/// A benchmark implementation describes: +/// - the algorithm (`meta()["algo"]`), +/// - a JSON-formatted specialization identifier (other keys in `meta()`), +/// - and the code that performs the timed measurement (`run()`). +/// +/// The executor uses this interface to: +/// - validate and sort benchmarks, +/// - construct per-benchmark state objects, +/// - run kernels and collect performance data, +/// - and emit structured JSON results. struct benchmark_interface { - /** - * \brief Returns a JSON object describing the benchmark. - * - * The returned JSON must include: - * - "algo": canonical algorithm name, - * - other keys describing the specialization. - * - * All benchmarks queued for one executor run must have the same "algo". - */ + /// Returns a JSON object describing the benchmark. + /// + /// The returned JSON must include: + /// - "algo": canonical algorithm name, + /// - other keys describing the specialization. + /// + /// All benchmarks queued for one executor run must have the same "algo". virtual json meta() const = 0; - /** - * \brief Executes the benchmark using the provided state. - * - * Implementations allocate input/output data, perform any required setup, - * and launch the algorithm under test. Timing, iteration control, and - * result reporting are handled through the supplied `state` object. - */ + /// Executes the benchmark using the provided state. + /// + /// Implementations allocate input/output data, perform any required setup, + /// and launch the algorithm under test. virtual void run(state& state) = 0; /// Virtual destructor for polymorphic cleanup. virtual ~benchmark_interface() = default; }; -/** - * \brief Generates and manages `SeedCount` seeds, from a single seed. - */ +/// Generates and manages `SeedCount` seeds, given a single input seed. template class seeds { @@ -3594,34 +3177,25 @@ class seeds std::array m_seeds; }; // class seeds -/** - * \brief Executes a suite of GPU benchmarks with configurable parameters. - * - * The executor class handles command-line parsing, benchmark queueing, - * execution, and logging of results in JSON format. Supports tuning - * GPU and benchmark parameters, including batch sizes, durations, and - * temperature limits. - */ +/// Executes a suite of GPU benchmarks with configurable parameters. +/// +/// The executor class handles command-line parsing, benchmark queueing, +/// execution, and logging of results in JSON format. Supports tuning +/// GPU and benchmark parameters, including batch sizes, durations, and +/// temperature limits. class executor { public: - /** - * \brief Constructs the executor and initializes parsing and logging. - * \param argc Argument count from main(). - * \param argv Argument values from main(). - * \param settings Optional benchmark-specific settings. - * \param flags Optional flags controlling executor behavior. - * \param stream Optional HIP stream to run the benchmarks on. - */ + /// Constructs the executor, and runs setup code. executor(int argc, char* argv[], primbench::settings settings = {}, detail::flags::FlagTag flags = flags::none, - hipStream_t stream = hipStreamDefault) + stream_t stream = default_stream) : m_settings(settings) , m_flags(flags) , m_stream(stream) - , m_own_stream(stream == hipStreamDefault) + , m_own_stream(stream == default_stream) , m_cli(argc, argv) { get_logger().save_program_start_time(); @@ -3629,9 +3203,9 @@ class executor parse(); // If user did not provide a stream, create a fast private one. - // We can't use hipStreamDefault, as it synchronizes with the host. + // We can't use default_stream, as it synchronizes with the host. if(m_own_stream) - PRIMBENCH_HIP_CHECK(hipStreamCreate(&m_stream)); + PRIMBENCH_CHECK(detail::stream_create(&m_stream)); m_stream_blocker = std::make_unique(m_stream, @@ -3641,51 +3215,37 @@ class executor ~executor() { if(m_own_stream && m_stream != nullptr) - PRIMBENCH_HIP_CHECK(hipStreamDestroy(m_stream)); + PRIMBENCH_CHECK(detail::stream_destroy(m_stream)); } - /** - * \brief Queue a benchmark for execution. - * \tparam Benchmark Type of benchmark to queue. - * \tparam Args Argument types for benchmark constructor. - * \param args Arguments to forward to the benchmark constructor. - * \return true to allow usage in global static initialization. - */ + /// \brief Queue a benchmark for execution. + /// + /// \tparam Benchmark Type of benchmark to queue. + /// \tparam Args Argument types for benchmark constructor. + /// \param args Arguments to forward to the benchmark constructor. + /// + /// \return true, which allows the function to be called in global scope. template static bool queue(Args&&... args) { - static_specializations.push_back(std::make_unique(std::forward(args)...)); + specializations.push_back(std::make_unique(std::forward(args)...)); return true; } - /** - * \brief Queue benchmarks using an autotune bulk creation function. - * \tparam BulkCreateFunction Callable that populates static_specializations. - * \param fn Function that creates benchmarks. - * \return true to allow usage in global static initialization. - */ + /// \brief Queue benchmarks using an autotune bulk creation function. + /// + /// \tparam BulkCreateFunction Callable that populates specializations. + /// \param fn Function that creates benchmarks. + /// + /// \return true, which allows the function to be called in global scope. template static bool queue_autotune(BulkCreateFunction&& fn) { - std::forward(fn)(static_specializations); + std::forward(fn)(specializations); return true; } - /** - * \brief Run all queued benchmarks and print progress/results. - * - * Performs the following: - * - Ensures run() is called only once per algorithm executor. - * - Sorts benchmarks to achieve a consistent order. - * - Validates that all benchmarks have the same algorithm name (`algo`). - * - Verifies that at least one benchmark is queued. - * - Ensures that all human-readable specialization names (`name`) are unique. - * - Computes output column widths and prints the benchmark header. - * - Executes each benchmark and prints progress/results. - * - Outputs a summary after all benchmarks are complete. - * - * \throws Exits the program with EXIT_FAILURE on validation errors. - */ + /// Prepares and runs all queued benchmark specializations. void run() { static bool run_called = false; @@ -3701,129 +3261,32 @@ class executor // Capture custom arguments after all have been registered. m_settings.custom_args = m_cli.get_all_custom_options(); - // Only keep filtered specializations, based on their name. - std::regex pattern(m_settings.filter); - static_specializations.erase(std::remove_if(static_specializations.begin(), - static_specializations.end(), - [&pattern](const auto& spec) { - return !std::regex_search( - spec.get()->meta().serialize_name(), - pattern); - }), - static_specializations.end()); - - // Sort to get a consistent order. - std::sort(static_specializations.begin(), - static_specializations.end(), - [](const auto& l, const auto& r) - { return l->meta().serialize_name() < r->meta().serialize_name(); }); - - size_t specialization_count = static_specializations.size(); - - if(specialization_count == 0) - { - std::cerr << "Error: At least one benchmark must be queued\n"; - if(!m_settings.filter.empty()) - { - std::cerr << "Hint: The currently used --filter '" << m_settings.filter - << "' is likely incorrect\n"; - } - exit(EXIT_FAILURE); - } - - std::string algorithm; - try - { - algorithm = static_specializations.front()->meta().get("algo"); - - // Validate that benchmarks have identical algo names. - for(const auto& bp : static_specializations) - { - auto bp_algo = bp->meta().get("algo"); - - if(bp_algo != algorithm) - { - std::cerr - << "Error: All specializations must have identical `algo` names, but '" - << bp_algo << "' and '" << algorithm << "' are different\n"; - exit(EXIT_FAILURE); - } - } - } - catch(const std::out_of_range& e) - { - std::cerr - << "Error: All benchmarks must return an \"algo\" key from their meta() method\n"; - exit(EXIT_FAILURE); - } - - get_logger().init(algorithm, specialization_count, m_settings, m_flags, get_amdsmi()); + filter_specializations(); + ensure_specializations_exist(); + sort_specializations(); - // Determine max specialization width, and validate that every name is unique. - m_spec_col_width = 0; - std::unordered_set seen_names; + std::string algorithm = get_common_algorithm(); - for(const auto& bp : static_specializations) - { - std::string name = bp->meta().serialize_name(); - size_t len = name.size(); - if(len > m_spec_col_width) - m_spec_col_width = len; + get_logger().init(algorithm, specializations.size(), m_settings, m_flags, get_monitor()); - if(!seen_names.insert(name).second) - { - std::cerr << "Error: Algorithm '" << algorithm - << "' has multiple specializations with the name '" << name << "'\n"; - exit(EXIT_FAILURE); - } - } + m_specialization_column_width = compute_max_specialization_width(algorithm); - m_family_col_width - = std::string("Index/").size() + std::to_string(specialization_count).size(); + constexpr size_t min_width = sizeof("9h 59m 59s") - 1; - if(m_settings.dry) - { - detail::progress::print_dry_header(algorithm, - m_spec_col_width, - m_family_col_width, - specialization_count); - } - else - { - detail::progress::print_header(algorithm, - m_spec_col_width, - m_family_col_width, - specialization_count, - m_settings.noise_timeout_secs); - } + m_index_column_width = std::max(min_width, + std::string("Index/").size() + + std::to_string(specializations.size()).size()); - // Run all benchmarks. - size_t family_index = 0; - for(auto& b_unique_ptr : static_specializations) - { - auto b = b_unique_ptr.get(); - auto meta = b->meta(); - auto algo = meta.get("algo"); + m_print_index = specializations.size() > 1; - if(m_settings.dry) - { - output_dry_specialization(algo, meta, family_index); - } - else - { - auto state = new_state(algo, meta, family_index); - b->run(state); - } + print_header(algorithm); - family_index++; - } + run_all_specializations(); get_logger().output_summary(); } - /** - * \brief Parses a command-line argument. - */ + /// Parses a command-line argument. template T get(std::string_view name, const T& default_val, std::string_view description) { @@ -3831,9 +3294,7 @@ class executor } private: - /** - * \brief Parse optional arguments. - */ + /// Parse optional arguments. void parse() { auto& cli = m_cli; @@ -3962,17 +3423,6 @@ class executor exit(EXIT_FAILURE); } - s.output_hip_device_properties_context - = cli.get("output-hip-device-properties-context", - s.output_hip_device_properties_context, - "Output a `hip_device_properties` object in the context object, " - "containing details about the GPU."); - - s.output_amdsmi_context = cli.get( - "output-amdsmi-context", - s.output_amdsmi_context, - "Output an `amdsmi` object in the context object, containing details about the GPU."); - s.output_batches = cli.get( "output-batches", s.output_batches, @@ -4000,34 +3450,176 @@ class executor } } - /** - * \brief Create a benchmark state object for execution. - * \param algo Algorithm name. - * \param meta Specialization metadata. - * \param family_index Index of benchmark in family. - * \return Configured state object for the benchmark. - */ - state new_state(std::string_view algo, json meta, size_t family_index) + /// Only keep filtered specializations, based on their name. + void filter_specializations() + { + std::regex pattern(m_settings.filter); + specializations.erase(std::remove_if(specializations.begin(), + specializations.end(), + [&pattern](const auto& spec) { + return !std::regex_search( + spec.get()->meta().serialize_name(), + pattern); + }), + specializations.end()); + } + + /// Ensures that at least one specialization is queued. + void ensure_specializations_exist() + { + if(specializations.size() == 0) + { + std::cerr << "Error: At least one benchmark must be queued\n"; + if(!m_settings.filter.empty()) + { + std::cerr << "Hint: The currently used --filter '" << m_settings.filter + << "' is likely incorrect\n"; + } + exit(EXIT_FAILURE); + } + } + + /// Sorts all registered specializations alphabetically by name. + /// + /// When specializations are added in global scope across multiple + /// translation units, the order of initialization is unspecified. + /// Calling this function ensures a consistent alphabetical order + /// regardless of initialization order. + void sort_specializations() + { + std::sort(specializations.begin(), + specializations.end(), + [](const auto& l, const auto& r) + { return l->meta().serialize_name() < r->meta().serialize_name(); }); + } + + /// Returns the algorithm name shared by all specializations. + /// + /// Validates that every specialization in `specializations` has + /// the same `"algo"` value in its meta data. Exits with an error + /// message if the `"algo"` key is missing or if algo values differ. + std::string get_common_algorithm() + { + std::string algorithm; + try + { + algorithm = specializations.front()->meta().get("algo"); + + for(const auto& bp : specializations) + { + auto bp_algo = bp->meta().get("algo"); + if(bp_algo != algorithm) + { + std::cerr + << "Error: All specializations must have identical `algo` names, but '" + << bp_algo << "' and '" << algorithm << "' are different\n"; + exit(EXIT_FAILURE); + } + } + } + catch(const std::out_of_range& e) + { + std::cerr + << "Error: All benchmarks must return an \"algo\" key from their meta() method\n"; + exit(EXIT_FAILURE); + } + + return algorithm; + } + + /// Computes the maximum column width of specialization names + /// and ensures all specialization names are unique. + size_t compute_max_specialization_width(std::string_view algorithm) + { + size_t max_width = sizeof("Specialization") - 1; + std::unordered_set seen_names; + + for(const auto& bp : specializations) + { + std::string name = bp->meta().serialize_name(); + size_t len = name.size(); + if(len > max_width) + max_width = len; + + if(!seen_names.insert(name).second) + { + std::cerr << "Error: Algorithm '" << algorithm + << "' has multiple specializations with the name '" << name << "'\n"; + exit(EXIT_FAILURE); + } + } + + return max_width; + } + + /// Prints a (dry) header. + void print_header(std::string_view algorithm) + { + if(m_settings.dry) + { + detail::progress::print_dry_header(algorithm, + m_specialization_column_width, + m_index_column_width, + specializations.size()); + } + else + { + detail::progress::print_header(algorithm, + m_specialization_column_width, + m_index_column_width, + specializations.size(), + m_settings.noise_timeout_secs); + } + } + + /// Runs all benchmark specializations. + void run_all_specializations() + { + size_t specialization_index = 0; + for(auto& b_unique_ptr : specializations) + { + auto b = b_unique_ptr.get(); + auto meta = b->meta(); + auto algo = meta.get("algo"); + + if(m_settings.dry) + { + output_dry_specialization(algo, meta, specialization_index); + } + else + { + auto state = new_state(algo, meta, specialization_index); + b->run(state); + } + + specialization_index++; + } + } + + /// Create a benchmark state object for execution. + state new_state(std::string_view algo, json meta, size_t specialization_index) { return state(algo, std::move(meta), - family_index, + specialization_index, + specializations.size(), m_stream, get_logger(), - get_amdsmi(), + get_monitor(), *m_stream_blocker, m_settings, m_flags, - m_spec_col_width, - m_family_col_width, + m_specialization_column_width, + m_index_column_width, + m_print_index, m_cache, m_warmer); } - /** - * \brief Outputs a single dry specialization. - */ - void output_dry_specialization(std::string_view algo, const json& meta, size_t family_index) + /// Outputs a single dry specialization. + void output_dry_specialization(std::string_view algo, + const json& meta, + size_t specialization_index) { std::string name = meta.serialize_name(); std::string serialized_meta = meta.serialize(); @@ -4048,11 +3640,12 @@ class executor detail::progress::print_dry_progress(name, algo, - family_index, - m_spec_col_width, - m_family_col_width); + specialization_index, + m_specialization_column_width, + m_index_column_width, + m_print_index); - get_logger().output_specialization(family_index, + get_logger().output_specialization(specialization_index, name, serialized_meta, kernels_per_batch, @@ -4066,50 +3659,47 @@ class executor end_temp, elapsed_host_secs, elapsed_gpu_secs, - noise_timeout, - get_amdsmi()); + noise_timeout); } - /** - * \brief Returns the logger singleton. - */ + /// Returns the logger singleton. detail::logger& get_logger() { return detail::logger::instance(); } - /** - * \brief Returns the amdsmi singleton. - */ - detail::amdsmi& get_amdsmi() + /// Returns the monitor singleton. + detail::monitor& get_monitor() { - return detail::amdsmi::instance(); + return detail::monitor::instance(); } - /** - * This vector is static, allowing queue_autotune() - * to register all specializations of autotuned benchmarks in one place. - */ - inline static std::vector> static_specializations; + using specializations_t = std::vector>; + + /// This vector is static, as benchmarks can be registered in global scope. + inline static specializations_t specializations; - settings m_settings; /**< CLI user settings */ + settings m_settings; ///< CLI user settings. - detail::flags::FlagTag m_flags; /**< Executor flags */ + detail::flags::FlagTag m_flags; ///< Executor flags. - hipStream_t m_stream; /**< HIP stream used for execution */ - bool m_own_stream; /** Whether primbench should create its own stream */ + stream_t m_stream; ///< Stream used for execution. + bool m_own_stream; ///< Whether primbench should create its own stream. - detail::cli m_cli; /**< Command-line argument parser */ + detail::cli m_cli; ///< Command-line argument parser. std::unique_ptr - m_stream_blocker; /**< Stream blocker to serialize output */ + m_stream_blocker; ///< Stream blocker to serialize output. + + size_t m_specialization_column_width; ///< Column width for specialization names. + + size_t m_index_column_width; ///< Column width for specialization index. + bool m_print_index; ///< Whether to print the index column. - size_t m_spec_col_width; /**< Column width for specialization names */ - size_t m_family_col_width; /**< Column width for family index */ + detail::cache_thrasher m_cache; ///< Cache clearing utility. - detail::cache_thrasher m_cache = detail::cache_thrasher(); /**< Cache clearing utility */ - detail::gpu_warmer m_warmer - = detail::gpu_warmer(m_settings, get_amdsmi()); /**< GPU warm-up utility */ + /// GPU warm-up utility. + detail::gpu_warmer m_warmer = detail::gpu_warmer(m_settings, get_monitor()); }; // class executor } // namespace primbench