CodeQL / Analyze (actions) (push) Canceled after 0s
CodeQL / Analyze (python) (push) Canceled after 0s
pre-commit / pre-commit (push) Canceled after 0s
Multi-Arch CI / setup (push) Canceled after 0s
Publish build_manylinux_x86_64 images / publish_build_manylinux_x86_64 (push) Canceled after 0s
Setup (Multi-Arch) / setup (push) Canceled after 0s
Multi-Arch CI / Linux::${{ fromJSON(needs.setup.outputs.linux_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubi10 image / publish_no_rocm_image_ubi10 (push) Canceled after 0s
Multi-Arch CI / Windows::${{ fromJSON(needs.setup.outputs.windows_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Multi-Arch CI / Manifest Diff (push) Canceled after 0s
TheRock Manifest Diff Report / Generate Manifest Diff Report (push) Canceled after 0s
Multi-Arch CI / CI Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_ocl_rt images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_rocgdb images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Unit Tests / Unit Tests :: ubuntu-24.04 (push) Canceled after 0s
Unit Tests / Unit Tests :: windows-2022 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_media (push) Canceled after 0s
Unit Tests / Unit Tests Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_openmpi (push) Canceled after 0s
4015 lines
136 KiB
C++
4015 lines
136 KiB
C++
// Copyright (c) 2025-2026 Advanced Micro Devices, Inc. All rights reserved.
|
|
//
|
|
// Permission is hereby granted, free of charge, to any person obtaining a copy
|
|
// of this software and associated documentation files (the "Software"), to deal
|
|
// in the Software without restriction, including without limitation the rights
|
|
// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
|
// copies of the Software, and to permit persons to whom the Software is
|
|
// furnished to do so, subject to the following conditions:
|
|
//
|
|
// The above copyright notice and this permission notice shall be included in
|
|
// all copies or substantial portions of the Software.
|
|
//
|
|
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
|
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
|
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
|
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
|
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
|
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
|
// THE SOFTWARE.
|
|
|
|
#pragma once
|
|
|
|
#ifdef _WIN32
|
|
// Disables the min() and max() macros from windows.h
|
|
#define NOMINMAX
|
|
|
|
#include <io.h>
|
|
#include <windows.h>
|
|
#else
|
|
#include <unistd.h>
|
|
#endif
|
|
|
|
#ifdef __HIP__
|
|
#include <hip/hip_fp16.h>
|
|
#include <hip/hip_runtime.h>
|
|
#elif defined(__CUDACC__)
|
|
#include <cuda/std/chrono>
|
|
#else
|
|
#error "Unsupported GPU platform"
|
|
#endif
|
|
|
|
#ifdef PRIMBENCH_NO_MONITORING
|
|
#define PRIMBENCH_HAS_MONITORING 0
|
|
#else
|
|
#define PRIMBENCH_HAS_MONITORING 1
|
|
#endif
|
|
|
|
#if PRIMBENCH_HAS_MONITORING
|
|
#ifdef __HIP__
|
|
#include <amd_smi/amdsmi.h>
|
|
#elif defined(__CUDACC__)
|
|
#include <nvml.h>
|
|
#endif
|
|
#endif
|
|
|
|
#include <array>
|
|
#include <cassert>
|
|
#include <chrono>
|
|
#include <cstring>
|
|
#include <fstream>
|
|
#include <functional>
|
|
#include <initializer_list>
|
|
#include <iomanip>
|
|
#include <iostream>
|
|
#include <map>
|
|
#include <memory>
|
|
#include <numeric>
|
|
#include <random>
|
|
#include <regex>
|
|
#include <thread>
|
|
#include <type_traits>
|
|
#include <unordered_set>
|
|
#include <variant>
|
|
|
|
/// 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
|
|
|
|
/// Registers a custom type name, used by `primbench::name<T>()`.
|
|
#define PRIMBENCH_REGISTER_TYPE(TYPE, NAME) \
|
|
namespace primbench::detail \
|
|
{ \
|
|
template<> \
|
|
struct type_name<TYPE> \
|
|
{ \
|
|
static inline const char* name = NAME; \
|
|
}; \
|
|
}
|
|
|
|
/// Asserts equality between input and expected. Works for scalar arithmetic types,
|
|
/// iterable containers, and brace-enclosed initializer lists. The expected value and
|
|
/// an optional tolerance (tol, defaults to 0.0) are captured via __VA_ARGS__ to handle
|
|
/// initializer list commas. Prints file:line, an error message to stderr, and exits
|
|
/// on mismatch.
|
|
#define PRIMBENCH_ASSERT(input, ...) \
|
|
primbench::detail::assert_equal(__FILE__, __LINE__, (input), __VA_ARGS__)
|
|
|
|
#ifdef __HIP__
|
|
/// Exits the program with an error message if the given HIP API call returns a failure status.
|
|
#define PRIMBENCH_CHECK(condition) \
|
|
do \
|
|
{ \
|
|
hipError_t status = condition; \
|
|
if(status != hipSuccess) \
|
|
{ \
|
|
std::cerr << __FILE__ << ":" << __LINE__ \
|
|
<< ": HIP error: " << hipGetErrorString(status) << "\n"; \
|
|
exit(status); \
|
|
} \
|
|
} \
|
|
while(0)
|
|
#else
|
|
/// Exits the program with an error message if the given CUDA API call returns a failure status.
|
|
#define PRIMBENCH_CHECK(condition) \
|
|
do \
|
|
{ \
|
|
cudaError_t status = condition; \
|
|
if(status != cudaSuccess) \
|
|
{ \
|
|
std::cerr << __FILE__ << ":" << __LINE__ \
|
|
<< ": CUDA error: " << cudaGetErrorString(status) << "\n"; \
|
|
exit(status); \
|
|
} \
|
|
} \
|
|
while(0)
|
|
#endif
|
|
|
|
#if defined(__HIP__) && PRIMBENCH_HAS_MONITORING
|
|
/// Exits the program with an error message if the given AMD SMI API call returns a failure status.
|
|
#define PRIMBENCH_AMDSMI_CHECK(condition) \
|
|
do \
|
|
{ \
|
|
amdsmi_status_t status = condition; \
|
|
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__) && PRIMBENCH_HAS_MONITORING
|
|
/// Exits the program with an error message if the given NVML API call returns a failure status.
|
|
#define PRIMBENCH_NVML_CHECK(condition) \
|
|
do \
|
|
{ \
|
|
nvmlReturn_t status = condition; \
|
|
if(status != NVML_SUCCESS) \
|
|
{ \
|
|
std::cerr << __FILE__ << ":" << __LINE__ \
|
|
<< ": NVML error: " << nvmlErrorString(status) << "\n"; \
|
|
exit(status); \
|
|
} \
|
|
} \
|
|
while(0)
|
|
#endif
|
|
|
|
namespace primbench
|
|
{
|
|
|
|
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<typename... Args>
|
|
void log(Args&&... args);
|
|
|
|
/// 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_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::string, bool, double, int, unsigned int, size_t>;
|
|
std::map<std::string, custom_arg_value>
|
|
custom_args; ///< Custom user-registered arguments with types.
|
|
|
|
}; // struct settings
|
|
|
|
namespace detail
|
|
{
|
|
|
|
/// Fallback for primbench::name().
|
|
template<class T>
|
|
struct type_name
|
|
{
|
|
static_assert(sizeof(T) == 0, "Call PRIMBENCH_REGISTER_TYPE() for this type");
|
|
|
|
static inline const char* name = "";
|
|
};
|
|
|
|
#ifdef _WIN32
|
|
/// Required to print horizontal_bar (u8"─").
|
|
inline const auto init_console_utf8 = []
|
|
{
|
|
SetConsoleOutputCP(CP_UTF8);
|
|
return 0;
|
|
}();
|
|
#endif
|
|
|
|
/// Returns whether the environment variable is set.
|
|
inline bool get_env(std::string_view key)
|
|
{
|
|
#ifdef _WIN32
|
|
char* value = nullptr;
|
|
size_t len = 0;
|
|
|
|
if(_dupenv_s(&value, &len, key.data()) != 0 || value == nullptr)
|
|
return false;
|
|
|
|
free(value);
|
|
return true;
|
|
#else
|
|
return std::getenv(key.data());
|
|
#endif
|
|
}
|
|
|
|
/// Caches whether ANSI color output is enabled.
|
|
/// The standard is described at https://bixense.com/clicolors/
|
|
inline bool use_color()
|
|
{
|
|
static const bool result = []
|
|
{
|
|
if(get_env("NO_COLOR"))
|
|
return false;
|
|
if(get_env("CLICOLOR_FORCE"))
|
|
return true;
|
|
|
|
#ifdef _WIN32
|
|
return bool(_isatty(_fileno(stdout)));
|
|
#else
|
|
return bool(isatty(fileno(stdout)));
|
|
#endif
|
|
}();
|
|
return result;
|
|
}
|
|
|
|
/// Clears the current line if colors are enabled.
|
|
inline std::ostream& clearline(std::ostream& os)
|
|
{
|
|
if(use_color())
|
|
os << "\r\033[K";
|
|
return os;
|
|
}
|
|
|
|
/// Resets the output text formatting to default if colors are enabled.
|
|
inline std::ostream& reset(std::ostream& os)
|
|
{
|
|
if(use_color())
|
|
os << "\033[0m";
|
|
return os;
|
|
}
|
|
|
|
/// Sets the output text color to gray if colors are enabled.
|
|
inline std::ostream& gray(std::ostream& os)
|
|
{
|
|
if(use_color())
|
|
os << "\033[90m";
|
|
return os;
|
|
}
|
|
|
|
/// Sets the output text color to green if colors are enabled.
|
|
inline std::ostream& green(std::ostream& os)
|
|
{
|
|
if(use_color())
|
|
os << "\033[92m";
|
|
return os;
|
|
}
|
|
|
|
/// Sets the output text color to red if colors are enabled.
|
|
inline std::ostream& red(std::ostream& os)
|
|
{
|
|
if(use_color())
|
|
os << "\033[91m";
|
|
return os;
|
|
}
|
|
|
|
/// Sets the output text color to yellow if colors are enabled.
|
|
inline std::ostream& yellow(std::ostream& os)
|
|
{
|
|
if(use_color())
|
|
os << "\033[93m";
|
|
return os;
|
|
}
|
|
|
|
/// Sets the output text color to blue if colors are enabled.
|
|
inline std::ostream& blue(std::ostream& os)
|
|
{
|
|
if(use_color())
|
|
os << "\033[94m";
|
|
return os;
|
|
}
|
|
|
|
template<typename T>
|
|
std::string value_to_string(const T& v)
|
|
{
|
|
return std::to_string(v);
|
|
}
|
|
|
|
inline std::string value_to_string(const __half& v)
|
|
{
|
|
return std::to_string(__half2float(v));
|
|
}
|
|
|
|
template<typename T>
|
|
constexpr bool is_numeric
|
|
= std::is_arithmetic_v<std::remove_cv_t<T>> || std::is_same_v<std::remove_cv_t<T>, __half>;
|
|
|
|
template<typename T>
|
|
constexpr bool is_fp_or_half
|
|
= std::is_floating_point_v<std::remove_cv_t<T>> || std::is_same_v<std::remove_cv_t<T>, __half>;
|
|
|
|
/// Asserts equality between two scalar values.
|
|
/// For floating-point types an optional tolerance can be specified.
|
|
/// Prints an error message to stderr and exits on mismatch.
|
|
template<typename T, typename U>
|
|
std::enable_if_t<is_numeric<T> && is_numeric<U>>
|
|
assert_equal(const char* file, int line, const T& input, const U& expected, double tol = 0.0)
|
|
{
|
|
clearline(std::cout);
|
|
const std::string prefix = std::string(file) + ":" + std::to_string(line) + ": ";
|
|
if constexpr(is_fp_or_half<T> || is_fp_or_half<U>)
|
|
{
|
|
const double diff = std::abs(static_cast<double>(input) - static_cast<double>(expected));
|
|
if(diff > tol)
|
|
{
|
|
std::cerr << prefix
|
|
<< "primbench::assert_equal() failed: Expected " + value_to_string(expected)
|
|
+ ", got " + value_to_string(input) + " (diff " + value_to_string(diff)
|
|
+ ", tol " + std::to_string(tol) + ")\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
}
|
|
else
|
|
{
|
|
if(input != expected)
|
|
{
|
|
std::cerr << prefix
|
|
<< "primbench::assert_equal() failed: Expected " + value_to_string(expected)
|
|
+ ", got " + value_to_string(input) + "\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
}
|
|
}
|
|
|
|
/// Asserts element-wise equality between two containers.
|
|
/// For floating-point types an optional tolerance can be specified.
|
|
/// Prints an error message to stderr and exits on size mismatch or value mismatch,
|
|
/// with a message indicating the index and the mismatching values.
|
|
template<typename ContainerT, typename ContainerU>
|
|
std::enable_if_t<!is_numeric<ContainerT> && !is_numeric<ContainerU>>
|
|
assert_equal(const char* file,
|
|
int line,
|
|
const ContainerT& input,
|
|
const ContainerU& expected,
|
|
double tol = 0.0)
|
|
{
|
|
clearline(std::cout);
|
|
const std::string prefix = std::string(file) + ":" + std::to_string(line) + ": ";
|
|
if(input.size() != expected.size())
|
|
{
|
|
std::cerr << prefix
|
|
<< "primbench::assert_equal() failed: Expected " + std::to_string(expected.size())
|
|
+ " values, got " + std::to_string(input.size()) + "\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
using T = typename ContainerT::value_type;
|
|
using U = typename ContainerU::value_type;
|
|
for(size_t i = 0; i < input.size(); ++i)
|
|
{
|
|
if constexpr(is_fp_or_half<T> || is_fp_or_half<U>)
|
|
{
|
|
const double diff
|
|
= std::abs(static_cast<double>(input[i]) - static_cast<double>(expected[i]));
|
|
if(diff > tol)
|
|
{
|
|
std::cerr << prefix
|
|
<< "primbench::assert_equal() failed at index " + std::to_string(i)
|
|
+ ": Expected " + value_to_string(expected[i]) + ", got "
|
|
+ value_to_string(input[i]) + " (diff " + value_to_string(diff)
|
|
+ ", tol " + std::to_string(tol) + ")\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
}
|
|
else
|
|
{
|
|
if(input[i] != expected[i])
|
|
{
|
|
std::cerr << prefix
|
|
<< "primbench::assert_equal() failed at index " + std::to_string(i)
|
|
+ ": Expected " + value_to_string(expected[i]) + ", got "
|
|
+ value_to_string(input[i]) + "\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
}
|
|
}
|
|
}
|
|
|
|
/// Asserts element-wise equality against an initializer list.
|
|
/// For floating-point types an optional tolerance can be specified.
|
|
/// Prints an error message to stderr and exits on size mismatch or value mismatch,
|
|
/// with a message indicating the index and the mismatching values.
|
|
template<typename ContainerT>
|
|
std::enable_if_t<!is_numeric<ContainerT>>
|
|
assert_equal(const char* file,
|
|
int line,
|
|
const ContainerT& input,
|
|
std::initializer_list<typename ContainerT::value_type> expected,
|
|
double tol = 0.0)
|
|
{
|
|
// Forward directly to the generic container overload.
|
|
assert_equal(file, line, input, std::vector(expected), tol);
|
|
}
|
|
|
|
/// 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);
|
|
}
|
|
|
|
template<typename F>
|
|
inline error_t occupancy_max_potential_block_size(int* min_grid_size,
|
|
int* block_size,
|
|
F kernel,
|
|
size_t dynamic_shared_memory_per_block,
|
|
int block_size_limit)
|
|
{
|
|
return PRIMBENCH_PREFIX_BACKEND(OccupancyMaxPotentialBlockSize)(min_grid_size,
|
|
block_size,
|
|
kernel,
|
|
dynamic_shared_memory_per_block,
|
|
block_size_limit);
|
|
}
|
|
|
|
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__
|
|
/// Serializes information about the used GPU.
|
|
inline std::string serialize_gpu_info()
|
|
{
|
|
std::ostringstream ss;
|
|
ss << "{";
|
|
|
|
int hip_device;
|
|
PRIMBENCH_CHECK(hipGetDevice(&hip_device));
|
|
|
|
hipDeviceProp_t dev_prop;
|
|
PRIMBENCH_CHECK(hipGetDeviceProperties(&dev_prop, hip_device));
|
|
ss << "\"name\":\"" << dev_prop.name << "\"";
|
|
|
|
std::string arch = dev_prop.gcnArchName;
|
|
size_t pos = arch.find(':');
|
|
if(pos != std::string::npos)
|
|
{
|
|
arch = arch.substr(0, pos);
|
|
}
|
|
ss << ",\"arch\":\"" << arch << "\"";
|
|
|
|
char pci_bus_id_arr[32];
|
|
PRIMBENCH_CHECK(hipDeviceGetPCIBusId(pci_bus_id_arr, sizeof(pci_bus_id_arr), hip_device));
|
|
ss << ",\"pci_bus_id\":\"" << pci_bus_id_arr << "\"";
|
|
|
|
ss << "}";
|
|
return ss.str();
|
|
}
|
|
#elif defined(__CUDACC__)
|
|
/// Serializes information about the used GPU.
|
|
inline std::string serialize_gpu_info()
|
|
{
|
|
std::ostringstream ss;
|
|
ss << "{";
|
|
|
|
int cuda_device;
|
|
PRIMBENCH_CHECK(cudaGetDevice(&cuda_device));
|
|
|
|
cudaDeviceProp dev_prop;
|
|
PRIMBENCH_CHECK(cudaGetDeviceProperties(&dev_prop, 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), cuda_device));
|
|
ss << ",\"pci_bus_id\":\"" << pci_bus_id_arr << "\"";
|
|
|
|
ss << "}";
|
|
return ss.str();
|
|
}
|
|
#endif
|
|
|
|
#if defined(__HIP__) && PRIMBENCH_HAS_MONITORING
|
|
|
|
/// Wrapper for AMD SMI (System Management Interface) 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;
|
|
|
|
/// Singleton accessor.
|
|
static monitor& instance()
|
|
{
|
|
static monitor instance;
|
|
return instance;
|
|
}
|
|
|
|
/// Saves the current temperature.
|
|
void save_start_temp()
|
|
{
|
|
m_saved_start_temp = get_temp();
|
|
}
|
|
|
|
/// Gets the saved starting temperature.
|
|
/// Specializations output this to JSON as `gpu_temp_celcius.start`.
|
|
uint16_t get_saved_start_temp() const
|
|
{
|
|
return m_saved_start_temp;
|
|
}
|
|
|
|
/// Returns the current temperature.
|
|
uint16_t get_temp() const
|
|
{
|
|
int64_t t;
|
|
PRIMBENCH_AMDSMI_CHECK(amdsmi_get_temp_metric(m_amdsmi_device,
|
|
get_temperature_type(),
|
|
AMDSMI_TEMP_CURRENT,
|
|
&t));
|
|
return t;
|
|
}
|
|
|
|
/// Serializes information about the monitoring library.
|
|
std::string serialize_monitoring() const
|
|
{
|
|
std::ostringstream ss;
|
|
ss << "{";
|
|
|
|
ss << "\"name\":\"amdsmi\"";
|
|
|
|
amdsmi_version_t amdsmi_version;
|
|
PRIMBENCH_AMDSMI_CHECK(amdsmi_get_lib_version(&amdsmi_version));
|
|
ss << ",\"version\":\"" << amdsmi_version.build << "\"";
|
|
|
|
ss << "}";
|
|
return ss.str();
|
|
}
|
|
|
|
/// Returns either "edge" or "hotspot".
|
|
std::string get_used_temperature_type_name() const
|
|
{
|
|
return get_temperature_type_name(get_temperature_type());
|
|
}
|
|
|
|
private:
|
|
monitor()
|
|
{
|
|
PRIMBENCH_AMDSMI_CHECK(amdsmi_init(AMDSMI_INIT_AMD_GPUS));
|
|
|
|
// 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();
|
|
}
|
|
|
|
~monitor()
|
|
{
|
|
PRIMBENCH_AMDSMI_CHECK(amdsmi_shut_down());
|
|
}
|
|
|
|
/// 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 temperature_type = [&]
|
|
{
|
|
const amdsmi_temperature_type_t types[] = {
|
|
AMDSMI_TEMPERATURE_TYPE_EDGE,
|
|
AMDSMI_TEMPERATURE_TYPE_HOTSPOT,
|
|
};
|
|
|
|
int64_t t;
|
|
for(auto type : types)
|
|
{
|
|
if(amdsmi_get_temp_metric(m_amdsmi_device, type, AMDSMI_TEMP_CURRENT, &t)
|
|
== AMDSMI_STATUS_SUCCESS)
|
|
{
|
|
return type;
|
|
}
|
|
}
|
|
|
|
std::cerr << "\nError: Failed to get any of these GPU temperatures: ";
|
|
bool first = true;
|
|
for(auto type : types)
|
|
{
|
|
if(!first)
|
|
std::cerr << ", ";
|
|
first = false;
|
|
std::cerr << get_temperature_type_name(type);
|
|
}
|
|
std::cerr << "\n";
|
|
exit(EXIT_FAILURE);
|
|
}();
|
|
|
|
return temperature_type;
|
|
}
|
|
|
|
/// Converts a temperature type enum to a string.
|
|
static const char* get_temperature_type_name(amdsmi_temperature_type_t type)
|
|
{
|
|
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);
|
|
}
|
|
}
|
|
|
|
/// Finds the AMD SMI processor handle matching the current HIP device.
|
|
amdsmi_processor_handle get_amdsmi_device() const
|
|
{
|
|
int hip_device;
|
|
PRIMBENCH_CHECK(hipGetDevice(&hip_device));
|
|
|
|
hipDeviceProp_t hip_props;
|
|
PRIMBENCH_CHECK(hipGetDeviceProperties(&hip_props, hip_device));
|
|
|
|
// Build the AMD SMI BDF struct from HIP device properties.
|
|
amdsmi_bdf_t addr{
|
|
.function_number = 0, // HIP doesn't expose PCI function ID.
|
|
.device_number = static_cast<uint8_t>(hip_props.pciDeviceID),
|
|
.bus_number = static_cast<uint8_t>(hip_props.pciBusID),
|
|
.domain_number = static_cast<uint16_t>(hip_props.pciDomainID),
|
|
};
|
|
|
|
amdsmi_processor_handle amdsmi_device;
|
|
PRIMBENCH_AMDSMI_CHECK(amdsmi_get_processor_handle_from_bdf(addr, &amdsmi_device));
|
|
|
|
return amdsmi_device;
|
|
}
|
|
|
|
/// Saved starting temperature for a specialization in state.run().
|
|
uint16_t m_saved_start_temp;
|
|
|
|
amdsmi_processor_handle m_amdsmi_device; ///< AMD SMI device.
|
|
}; // class monitor
|
|
|
|
#elif defined(__CUDACC__) && PRIMBENCH_HAS_MONITORING
|
|
|
|
/// Wrapper for 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;
|
|
|
|
/// Singleton accessor.
|
|
static monitor& instance()
|
|
{
|
|
static monitor instance;
|
|
return instance;
|
|
}
|
|
|
|
/// Saves the current temperature.
|
|
void save_start_temp()
|
|
{
|
|
m_saved_start_temp = get_temp();
|
|
}
|
|
|
|
/// Gets the saved starting temperature.
|
|
/// Specializations output this to JSON as `gpu_temp_celcius.start`.
|
|
uint16_t get_saved_start_temp() const
|
|
{
|
|
return m_saved_start_temp;
|
|
}
|
|
|
|
/// Returns the current temperature.
|
|
uint16_t get_temp() const
|
|
{
|
|
unsigned int t;
|
|
PRIMBENCH_NVML_CHECK(nvmlDeviceGetTemperature(m_nvml_device, NVML_TEMPERATURE_GPU, &t));
|
|
return t;
|
|
}
|
|
|
|
/// Serializes information about the monitoring library.
|
|
std::string serialize_monitoring() const
|
|
{
|
|
std::ostringstream ss;
|
|
ss << "{";
|
|
|
|
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 << "}";
|
|
return ss.str();
|
|
}
|
|
|
|
/// NVML only exposes a single temperature sensor.
|
|
std::string get_used_temperature_type_name() const
|
|
{
|
|
return "gpu";
|
|
}
|
|
|
|
private:
|
|
monitor()
|
|
{
|
|
PRIMBENCH_NVML_CHECK(nvmlInit());
|
|
|
|
PRIMBENCH_NVML_CHECK(nvmlDeviceGetHandleByIndex(0, &m_nvml_device));
|
|
}
|
|
|
|
~monitor()
|
|
{
|
|
PRIMBENCH_NVML_CHECK(nvmlShutdown());
|
|
}
|
|
|
|
/// Saved starting temperature for a specialization in state.run().
|
|
uint16_t m_saved_start_temp;
|
|
|
|
nvmlDevice_t m_nvml_device; ///< NVML device.
|
|
}; // class monitor
|
|
|
|
#endif
|
|
|
|
/// Namespace for flag definitions and utilities.
|
|
namespace flags
|
|
{
|
|
|
|
/// Enum representing different flags.
|
|
enum class Flags : uint32_t
|
|
{
|
|
none = 0x0, ///< No flags set.
|
|
sync = 0x1, ///< Synchronization flag.
|
|
};
|
|
|
|
/// Wrapper for Flags with utility operations.
|
|
struct FlagTag
|
|
{
|
|
Flags value{Flags::none}; ///< Underlying flag value.
|
|
|
|
/// Construct from a specific flag.
|
|
constexpr FlagTag(Flags v) : value(v) {}
|
|
|
|
/// Bitwise OR operator for combining flags.
|
|
friend constexpr FlagTag operator|(FlagTag a, FlagTag b)
|
|
{
|
|
return FlagTag(
|
|
static_cast<Flags>(static_cast<uint32_t>(a.value) | static_cast<uint32_t>(b.value)));
|
|
}
|
|
|
|
/// Check if a flag is set.
|
|
constexpr bool has(FlagTag f) const
|
|
{
|
|
return (static_cast<uint32_t>(value) & static_cast<uint32_t>(f.value)) != 0;
|
|
}
|
|
};
|
|
|
|
} // namespace flags
|
|
|
|
/// 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.
|
|
logger(const logger&) = delete;
|
|
logger& operator=(const logger&) = delete;
|
|
logger(logger&&) = delete;
|
|
logger& operator=(logger&&) = delete;
|
|
|
|
/// Singleton accessor.
|
|
static logger& instance()
|
|
{
|
|
static logger instance;
|
|
return instance;
|
|
}
|
|
|
|
/// Saves the program start time.
|
|
void save_program_start_time()
|
|
{
|
|
m_program_start_time = std::chrono::steady_clock::now();
|
|
}
|
|
|
|
/// Returns the program start time.
|
|
std::chrono::time_point<std::chrono::steady_clock> 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)
|
|
{
|
|
m_output_batches = settings.output_batches;
|
|
m_spaces_per_indent = settings.spaces_per_indent;
|
|
m_outputting_csv = !settings.csv_out.empty();
|
|
|
|
m_json_out.open(std::string(settings.json_out), std::ios::out | std::ios::trunc);
|
|
if(!m_json_out)
|
|
{
|
|
std::cerr << "Error: Failed to open " << settings.json_out << " for writing\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
|
|
if(m_outputting_csv)
|
|
{
|
|
m_csv_out.open(std::string(settings.csv_out), std::ios::out | std::ios::trunc);
|
|
if(!m_csv_out)
|
|
{
|
|
std::cerr << "Error: Failed to open " << settings.csv_out << " for writing\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
}
|
|
|
|
m_json_out << indent(
|
|
serialize_json_prologue(algorithm, specialization_count, settings, flags),
|
|
0);
|
|
m_json_out << "[";
|
|
if(m_spaces_per_indent > 0)
|
|
m_json_out << "\n";
|
|
m_json_out.flush();
|
|
|
|
if(m_outputting_csv)
|
|
{
|
|
// Output the CSV header.
|
|
m_csv_out << "index,name,bytes_per_second,gib_per_second,items_per_second"
|
|
",noise_timeout,noise_percent\n";
|
|
m_csv_out.flush();
|
|
}
|
|
|
|
m_first_specialization = true;
|
|
}
|
|
|
|
/// Stores a batch of benchmark results.
|
|
void save(double batch_ms, const std::vector<float>& iterations_ms)
|
|
{
|
|
struct batch batch
|
|
{};
|
|
|
|
batch.batch_ms = batch_ms;
|
|
batch.iterations_ms = iterations_ms;
|
|
|
|
m_batches.push_back(batch);
|
|
}
|
|
|
|
/// Outputs JSON and CSV specialization information.
|
|
void output_specialization(size_t index,
|
|
std::string_view name,
|
|
std::string_view serialized_meta,
|
|
size_t kernels_per_batch,
|
|
double ms_per_batch,
|
|
double bytes_per_sec,
|
|
double items_per_sec,
|
|
size_t bytes_per_item,
|
|
size_t items,
|
|
double noise_percent,
|
|
double elapsed_host_secs,
|
|
double elapsed_gpu_secs,
|
|
bool noise_timeout)
|
|
{
|
|
output_json_specialization(index,
|
|
name,
|
|
serialized_meta,
|
|
kernels_per_batch,
|
|
ms_per_batch,
|
|
bytes_per_sec,
|
|
items_per_sec,
|
|
bytes_per_item,
|
|
items,
|
|
noise_percent,
|
|
elapsed_host_secs,
|
|
elapsed_gpu_secs,
|
|
noise_timeout);
|
|
|
|
if(m_outputting_csv)
|
|
{
|
|
output_csv_specialization(index,
|
|
name,
|
|
bytes_per_sec,
|
|
items_per_sec,
|
|
noise_timeout,
|
|
noise_percent);
|
|
}
|
|
|
|
m_total_elapsed_gpu_secs += elapsed_gpu_secs;
|
|
|
|
if(noise_timeout)
|
|
m_noise_timeouts++;
|
|
|
|
m_batches.clear();
|
|
}
|
|
|
|
/// Outputs a summary object at the end of the JSON file.
|
|
void output_summary()
|
|
{
|
|
// Close the JSON file's outer array.
|
|
if(m_spaces_per_indent > 0)
|
|
{
|
|
m_json_out << "\n";
|
|
m_json_out << std::string(m_spaces_per_indent, ' ');
|
|
}
|
|
m_json_out << "],";
|
|
|
|
if(m_spaces_per_indent > 0)
|
|
m_json_out << "\n";
|
|
m_json_out << indent(serialize_summary(), 1);
|
|
|
|
// Close the JSON file's outer object.
|
|
if(m_spaces_per_indent > 0)
|
|
m_json_out << "\n";
|
|
m_json_out << "}";
|
|
if(m_spaces_per_indent > 0)
|
|
m_json_out << "\n";
|
|
|
|
m_json_out.close();
|
|
|
|
if(m_outputting_csv)
|
|
{
|
|
m_csv_out.close();
|
|
}
|
|
}
|
|
|
|
private:
|
|
logger() = default;
|
|
|
|
struct batch
|
|
{
|
|
double batch_ms; ///< Total time for the batch.
|
|
std::vector<float> iterations_ms; ///< Time per iteration.
|
|
};
|
|
|
|
/// 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)
|
|
{
|
|
std::ostringstream ss;
|
|
ss << "{";
|
|
|
|
ss << "\"context\":" << serialize_context(algorithm, specialization_count, settings, flags);
|
|
|
|
ss << ",";
|
|
if(m_spaces_per_indent > 0)
|
|
ss << "\n";
|
|
|
|
ss << "\"specializations\":";
|
|
|
|
return ss.str();
|
|
}
|
|
|
|
/// 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
|
|
{
|
|
std::ostringstream ss;
|
|
ss << "{";
|
|
|
|
ss << "\"results_version\":\"4.0.0\"";
|
|
ss << ",\"general\":" << serialize_general(algorithm, specialization_count);
|
|
ss << ",\"settings\":" << serialize_settings(settings);
|
|
|
|
std::string custom_cli = serialize_custom_settings(settings);
|
|
if(!custom_cli.empty())
|
|
{
|
|
ss << ",\"custom_settings\":" << custom_cli;
|
|
}
|
|
|
|
ss << ",\"flags\":" << serialize_flags(flags);
|
|
|
|
ss << "}";
|
|
return ss.str();
|
|
}
|
|
|
|
/// Serializes general benchmark context info into JSON.
|
|
std::string serialize_general(std::string_view algorithm, size_t specialization_count) const
|
|
{
|
|
std::ostringstream ss;
|
|
ss << "{";
|
|
|
|
ss << "\"algorithm\":\"" << algorithm << "\"";
|
|
ss << ",\"specialization_count\":" << specialization_count;
|
|
|
|
#ifdef NDEBUG
|
|
const char build_type[] = "release";
|
|
#else
|
|
const char build_type[] = "debug";
|
|
#endif
|
|
ss << ",\"library_build_type\":\"" << build_type << "\"";
|
|
|
|
ss << ",\"gpu\":" << serialize_gpu_info();
|
|
|
|
ss << ",\"backend\":" << serialize_backend();
|
|
|
|
#if PRIMBENCH_HAS_MONITORING
|
|
ss << ",\"monitoring\":" << monitor::instance().serialize_monitoring();
|
|
|
|
ss << ",\"temperature_type\":\"" << monitor::instance().get_used_temperature_type_name()
|
|
<< "\"";
|
|
#endif
|
|
|
|
ss << ",\"host_name\":\"" << get_host_name() << "\"";
|
|
|
|
ss << ",\"date\":\"" << date() << "\"";
|
|
|
|
#ifdef BRANCH_NAME
|
|
ss << ",\"branch_name\":\"" << BRANCH_NAME << "\"";
|
|
#endif
|
|
#ifdef COMMIT_HASH
|
|
ss << ",\"commit_hash\":\"" << COMMIT_HASH << "\"";
|
|
#endif
|
|
|
|
ss << "}";
|
|
return ss.str();
|
|
}
|
|
|
|
#ifdef __HIP__
|
|
/// Serializes information about the HIP backend.
|
|
std::string serialize_backend() const
|
|
{
|
|
std::ostringstream ss;
|
|
ss << "{";
|
|
|
|
// Backend name.
|
|
ss << "\"name\":\"hip\"";
|
|
|
|
// HIP version.
|
|
ss << ",\"hip_version\":\"" << HIP_VERSION_MAJOR << "." << HIP_VERSION_MINOR << "."
|
|
<< HIP_VERSION_PATCH << "-" << HIP_VERSION_GITHASH << "\"";
|
|
|
|
// 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 << "}";
|
|
return ss.str();
|
|
}
|
|
#elif defined(__CUDACC__)
|
|
/// Serializes information about the CUDA backend.
|
|
std::string serialize_backend() const
|
|
{
|
|
std::ostringstream ss;
|
|
ss << "{";
|
|
|
|
// 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 << "}";
|
|
|
|
ss << "}";
|
|
return ss.str();
|
|
}
|
|
#endif
|
|
|
|
static std::string get_host_name()
|
|
{
|
|
#ifdef _WIN32
|
|
char buffer[MAX_COMPUTERNAME_LENGTH + 1] = {0};
|
|
DWORD size = sizeof(buffer);
|
|
if(!GetComputerNameA(buffer, &size))
|
|
{
|
|
std::cerr << "Error: GetComputerNameA failed with Win32 Error Code: " << GetLastError()
|
|
<< "\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
#else
|
|
char buffer[HOST_NAME_MAX + 1] = {0};
|
|
// POSIX allows truncation without a null terminator.
|
|
if(gethostname(buffer, HOST_NAME_MAX) != 0)
|
|
{
|
|
std::cerr << "Error: gethostname failed: " << std::strerror(errno) << "\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
#endif
|
|
return std::string(buffer);
|
|
}
|
|
|
|
/// Returns the local date and time as an RFC3339 string: `yyyy-mm-ddTHH:MM:SS±HH:MM`.
|
|
static std::string date()
|
|
{
|
|
using namespace std::chrono;
|
|
|
|
// 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.
|
|
std::tm local_tm{};
|
|
#ifdef _WIN32
|
|
localtime_s(&local_tm, &now_c);
|
|
#else
|
|
localtime_r(&now_c, &local_tm);
|
|
#endif
|
|
|
|
// Format date and time.
|
|
std::ostringstream oss;
|
|
oss << std::put_time(&local_tm, "%Y-%m-%dT%H:%M:%S");
|
|
|
|
// Compute timezone offset.
|
|
std::tm utc_tm{};
|
|
#ifdef _WIN32
|
|
gmtime_s(&utc_tm, &now_c);
|
|
#else
|
|
gmtime_r(&now_c, &utc_tm);
|
|
#endif
|
|
|
|
// Offset in seconds.
|
|
int offset_sec
|
|
= static_cast<int>(std::difftime(std::mktime(&local_tm), std::mktime(&utc_tm)));
|
|
char sign = offset_sec >= 0 ? '+' : '-';
|
|
offset_sec = std::abs(offset_sec);
|
|
int offset_hours = offset_sec / 3600;
|
|
int offset_minutes = (offset_sec % 3600) / 60;
|
|
|
|
// Append timezone.
|
|
oss << sign << std::setw(2) << std::setfill('0') << offset_hours << ':' << std::setw(2)
|
|
<< std::setfill('0') << offset_minutes;
|
|
|
|
return oss.str();
|
|
}
|
|
|
|
/// Serializes CLI settings into JSON.
|
|
std::string serialize_settings(const settings& settings) const
|
|
{
|
|
std::ostringstream ss;
|
|
ss << "{";
|
|
|
|
const auto& s = settings;
|
|
|
|
ss << "\"size\":" << s.size;
|
|
ss << ",\"hot\":" << std::boolalpha << s.hot;
|
|
ss << ",\"seed\":" << s.seed;
|
|
ss << ",\"json_out\":\"" << s.json_out << "\"";
|
|
ss << ",\"csv_out\":\"" << s.csv_out << "\"";
|
|
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;
|
|
ss << ",\"noise_timeout_secs\":" << s.noise_timeout_secs;
|
|
ss << ",\"batch_window_size\":" << s.batch_window_size;
|
|
ss << ",\"noise_tolerance_percent\":" << s.noise_tolerance_percent;
|
|
ss << ",\"min_gpu_temp\":" << s.min_gpu_temp;
|
|
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_batches\":" << s.output_batches;
|
|
ss << ",\"spaces_per_indent\":" << s.spaces_per_indent;
|
|
ss << ",\"stream_blocking_timeout_secs\":" << s.stream_blocking_timeout_secs;
|
|
|
|
ss << "}";
|
|
return ss.str();
|
|
}
|
|
|
|
/// 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())
|
|
{
|
|
return "";
|
|
}
|
|
|
|
std::ostringstream ss;
|
|
ss << "{";
|
|
ss << std::boolalpha;
|
|
|
|
bool first = true;
|
|
for(const auto& [key, value] : settings.custom_args)
|
|
{
|
|
if(!first)
|
|
ss << ",";
|
|
first = false;
|
|
|
|
ss << "\"" << key << "\":";
|
|
|
|
std::visit(
|
|
[&](auto const& val)
|
|
{
|
|
using V = std::decay_t<decltype(val)>;
|
|
if constexpr(std::is_same_v<V, std::string>)
|
|
{
|
|
ss << "\"" << val << "\"";
|
|
}
|
|
else
|
|
{
|
|
ss << val;
|
|
}
|
|
},
|
|
value);
|
|
}
|
|
|
|
ss << "}";
|
|
return ss.str();
|
|
}
|
|
|
|
/// Formats a JSON string with indentation.
|
|
std::string indent(std::string_view str, size_t starting_indent_level)
|
|
{
|
|
if(m_spaces_per_indent == 0)
|
|
return std::string(str);
|
|
|
|
std::ostringstream out;
|
|
size_t indent_level = starting_indent_level;
|
|
bool in_string = false;
|
|
|
|
if(indent_level > 0)
|
|
out << std::string(indent_level * m_spaces_per_indent, ' ');
|
|
|
|
for(size_t i = 0; i < str.size(); ++i)
|
|
{
|
|
char c = str[i];
|
|
|
|
if(c == '\"')
|
|
{
|
|
// Detect escaped quotes.
|
|
bool escaped = false;
|
|
size_t j = i;
|
|
while(j > 0 && str[--j] == '\\')
|
|
escaped = !escaped;
|
|
if(!escaped)
|
|
in_string = !in_string;
|
|
out << c;
|
|
}
|
|
else if(!in_string)
|
|
{
|
|
switch(c)
|
|
{
|
|
case '{':
|
|
case '[':
|
|
out << c << '\n';
|
|
++indent_level;
|
|
out << std::string(indent_level * m_spaces_per_indent, ' ');
|
|
break;
|
|
case '}':
|
|
case ']':
|
|
out << '\n';
|
|
if(indent_level > 0)
|
|
--indent_level;
|
|
out << std::string(indent_level * m_spaces_per_indent, ' ') << c;
|
|
break;
|
|
case ',':
|
|
out << c << '\n' << std::string(indent_level * m_spaces_per_indent, ' ');
|
|
break;
|
|
case ':': out << c << ' '; break;
|
|
default:
|
|
if(!isspace(static_cast<unsigned char>(c)))
|
|
out << c;
|
|
break;
|
|
}
|
|
}
|
|
else
|
|
{
|
|
out << c;
|
|
}
|
|
}
|
|
|
|
return out.str();
|
|
}
|
|
|
|
/// Serializes benchmark flags into JSON.
|
|
std::string serialize_flags(flags::FlagTag flags) const
|
|
{
|
|
std::ostringstream ss;
|
|
ss << "{";
|
|
ss << "\"sync\":" << std::boolalpha << flags.has(flags::Flags::sync);
|
|
ss << "}";
|
|
return ss.str();
|
|
}
|
|
|
|
/// Outputs specialization information to JSON.
|
|
void output_json_specialization(size_t index,
|
|
std::string_view name,
|
|
std::string_view serialized_meta,
|
|
size_t kernels_per_batch,
|
|
double ms_per_batch,
|
|
double bytes_per_sec,
|
|
double items_per_sec,
|
|
size_t bytes_per_item,
|
|
size_t items,
|
|
double noise_percent,
|
|
double elapsed_host_secs,
|
|
double elapsed_gpu_secs,
|
|
bool noise_timeout)
|
|
{
|
|
// Specializations need a comma between them.
|
|
if(!m_first_specialization)
|
|
{
|
|
m_json_out << ",";
|
|
if(m_spaces_per_indent > 0)
|
|
m_json_out << "\n";
|
|
}
|
|
else
|
|
m_first_specialization = false;
|
|
|
|
m_json_out << indent(serialize_specialization(index,
|
|
name,
|
|
serialized_meta,
|
|
kernels_per_batch,
|
|
ms_per_batch,
|
|
bytes_per_sec,
|
|
items_per_sec,
|
|
bytes_per_item,
|
|
items,
|
|
noise_percent,
|
|
elapsed_host_secs,
|
|
elapsed_gpu_secs,
|
|
noise_timeout),
|
|
2);
|
|
|
|
m_json_out.flush();
|
|
}
|
|
|
|
/// Outputs specialization information to CSV.
|
|
void output_csv_specialization(size_t index,
|
|
std::string_view name,
|
|
double bytes_per_sec,
|
|
double items_per_sec,
|
|
bool noise_timeout,
|
|
double noise_percent)
|
|
{
|
|
m_csv_out << index << ",\"" << name << "\"" << "," << bytes_per_sec << ","
|
|
<< bytes_per_sec / (1024.0 * 1024.0 * 1024.0) << "," << items_per_sec << ","
|
|
<< noise_timeout << "," << noise_percent << "\n"
|
|
<< std::flush;
|
|
}
|
|
|
|
/// Serializes a specialization into JSON format.
|
|
std::string serialize_specialization(size_t index,
|
|
std::string_view name,
|
|
std::string_view serialized_meta,
|
|
size_t kernels_per_batch,
|
|
double ms_per_batch,
|
|
double bytes_per_sec,
|
|
double items_per_sec,
|
|
size_t bytes_per_item,
|
|
size_t items,
|
|
double noise_percent,
|
|
double elapsed_host_secs,
|
|
double elapsed_gpu_secs,
|
|
bool noise_timeout) const
|
|
{
|
|
std::ostringstream ss;
|
|
ss << "{";
|
|
|
|
ss << "\"index\":" << index;
|
|
ss << ",\"name\":\"" << name << "\"";
|
|
|
|
ss << ",\"bytes_per_second\":" << bytes_per_sec;
|
|
ss << ",\"items_per_second\":" << items_per_sec;
|
|
|
|
ss << ",\"bytes_per_item\":" << bytes_per_item;
|
|
ss << ",\"items\":" << items;
|
|
|
|
ss << ",\"noise_timeout\":" << std::boolalpha << noise_timeout;
|
|
ss << ",\"noise_percent\":" << noise_percent;
|
|
|
|
ss << ",\"meta\":" << serialized_meta;
|
|
|
|
ss << ",\"elapsed_secs\":{";
|
|
ss << "\"host\":" << elapsed_host_secs;
|
|
ss << ",\"gpu\":" << elapsed_gpu_secs;
|
|
ss << "}";
|
|
|
|
#if PRIMBENCH_HAS_MONITORING
|
|
ss << ",\"gpu_temp_celsius\":{";
|
|
ss << "\"start\":" << monitor::instance().get_saved_start_temp();
|
|
ss << ",\"end\":" << monitor::instance().get_temp();
|
|
ss << "}";
|
|
#endif
|
|
|
|
ss << ",\"calls\":{";
|
|
ss << "\"kernel_calls_per_batch\":" << kernels_per_batch;
|
|
ss << ",\"ms_per_batch\":" << ms_per_batch;
|
|
ss << ",\"batches\":" << m_batches.size();
|
|
ss << ",\"kernel_calls\":" << kernels_per_batch * m_batches.size();
|
|
ss << "}";
|
|
|
|
if(m_output_batches)
|
|
{
|
|
ss << ",\"batches\":[";
|
|
for(size_t i = 0; i < m_batches.size(); ++i)
|
|
{
|
|
if(i > 0)
|
|
ss << ",";
|
|
const auto& batch = m_batches[i];
|
|
ss << "{";
|
|
ss << "\"batch_ms\":" << batch.batch_ms;
|
|
ss << ",\"iterations_ms\":" << serialize_iterations_ms(batch.iterations_ms);
|
|
ss << "}";
|
|
}
|
|
ss << "]";
|
|
}
|
|
|
|
ss << "}";
|
|
return ss.str();
|
|
}
|
|
|
|
/// Serializes iteration times into JSON array.
|
|
std::string serialize_iterations_ms(const std::vector<float>& iterations_ms) const
|
|
{
|
|
std::ostringstream ss;
|
|
ss << "[";
|
|
for(size_t i = 0; i < iterations_ms.size(); ++i)
|
|
{
|
|
if(i > 0)
|
|
ss << ",";
|
|
ss << iterations_ms[i];
|
|
}
|
|
ss << "]";
|
|
return ss.str();
|
|
}
|
|
|
|
/// Serializes summary info into JSON format.
|
|
std::string serialize_summary() const
|
|
{
|
|
std::ostringstream ss;
|
|
ss << "\"summary\":{";
|
|
|
|
ss << "\"noise_timeouts\":" << m_noise_timeouts;
|
|
|
|
auto now = std::chrono::steady_clock::now();
|
|
auto elapsed_host = now - m_program_start_time;
|
|
|
|
double elapsed_host_secs = std::chrono::duration<double>(elapsed_host).count();
|
|
|
|
ss << ",\"elapsed_secs\":{";
|
|
ss << "\"host\":" << elapsed_host_secs;
|
|
ss << ",\"gpu\":" << m_total_elapsed_gpu_secs;
|
|
ss << "}";
|
|
|
|
ss << "}";
|
|
return ss.str();
|
|
}
|
|
|
|
std::chrono::time_point<std::chrono::steady_clock> 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<batch> 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
|
|
|
|
/// Provides functions for progress display and formatting during GPU benchmarking.
|
|
namespace progress
|
|
{
|
|
// 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 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_column_width = status_header.size();
|
|
|
|
std::cout << std::setw(status_column_width) << std::left << status_header << " "
|
|
<< std::setw(specialization_column_width) << std::left << "Specialization";
|
|
|
|
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";
|
|
}
|
|
|
|
/// Prints the table header for algorithm progress output.
|
|
inline void print_header(std::string_view algo_name,
|
|
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_column_width = std::max(status_header.size(), noisy_status.size());
|
|
|
|
std::cout << std::setw(status_column_width) << std::left << status_header << " "
|
|
<< std::setw(noise_column_width) << std::left << "Noise";
|
|
|
|
#if PRIMBENCH_HAS_MONITORING
|
|
std::cout << " " << std::setw(gpu_temp_column_width) << std::left << "GPU °C";
|
|
#endif
|
|
|
|
std::cout << " " << std::setw(bytes_per_sec_column_width) << std::left << "Bytes/sec" << " "
|
|
<< std::setw(specialization_column_width) << std::left << "Specialization";
|
|
|
|
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;
|
|
underline_width += 2 + noise_column_width;
|
|
#if PRIMBENCH_HAS_MONITORING
|
|
underline_width += 2 + gpu_temp_column_width;
|
|
#endif
|
|
underline_width += 2 + bytes_per_sec_column_width;
|
|
underline_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";
|
|
}
|
|
|
|
/// 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;
|
|
}
|
|
|
|
/// 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;
|
|
}
|
|
|
|
/// 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 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_column_width = status_header.size();
|
|
|
|
line << clearline << green << std::setw(status_column_width) << std::left << "Success" << reset;
|
|
|
|
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;
|
|
}
|
|
|
|
/// 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<int>(remaining_secs) / 3600;
|
|
int m = (static_cast<int>(remaining_secs) % 3600) / 60;
|
|
int s = static_cast<int>(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,
|
|
std::string_view status_msg,
|
|
std::string_view specialization,
|
|
std::string_view algo_name,
|
|
uint64_t batch_window_size,
|
|
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,
|
|
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_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_column_width)
|
|
{
|
|
bar_width = status_column_width - batch_str.size() - 1;
|
|
}
|
|
|
|
std::ostringstream line;
|
|
line << clearline << batch_str;
|
|
|
|
// Progress bar (only shown during iteration).
|
|
if(status_msg.empty())
|
|
{
|
|
line << " ";
|
|
|
|
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.
|
|
size_t yellow_chars = 0;
|
|
if(filled >= bar_width)
|
|
{
|
|
double frac = std::min(1.0, elapsed_host_secs / noise_timeout_secs);
|
|
yellow_chars = bar_width * frac;
|
|
}
|
|
|
|
// Build the progress bar in one pass.
|
|
for(size_t j = 0; j < bar_width; ++j)
|
|
{
|
|
if(j < yellow_chars)
|
|
line << yellow << horizontal_bar;
|
|
else if(j < filled)
|
|
line << green << horizontal_bar;
|
|
else
|
|
line << gray << horizontal_bar;
|
|
}
|
|
line << reset;
|
|
}
|
|
|
|
// Alignment for subsequent columns.
|
|
size_t used = batch_str.size() + status_msg.empty() + (status_msg.empty() ? bar_width : 0);
|
|
if(used < status_column_width)
|
|
line << std::string(status_column_width - used, ' ');
|
|
|
|
// Noise %.
|
|
std::ostringstream percent_stream;
|
|
|
|
if(noise_percent >= 100.0)
|
|
{
|
|
percent_stream << static_cast<int>(noise_percent);
|
|
}
|
|
else
|
|
{
|
|
percent_stream << std::fixed << std::setprecision(1) << noise_percent;
|
|
}
|
|
|
|
auto& noise_color = (noise_percent > noise_tolerance_percent) ? red : reset;
|
|
|
|
line << " " << noise_color << std::setw(noise_column_width - sizeof('%')) << std::right
|
|
<< percent_stream.str() << "%" << reset;
|
|
|
|
#if PRIMBENCH_HAS_MONITORING
|
|
// GPU temperature.
|
|
line << " " << std::setw(gpu_temp_column_width) << std::right
|
|
<< monitor::instance().get_temp();
|
|
#endif
|
|
|
|
// Bytes/sec.
|
|
line << " " << std::setw(bytes_per_sec_column_width) << std::right << std::scientific
|
|
<< std::setprecision(2) << bytes_per_sec;
|
|
|
|
// Specialization.
|
|
line << " " << std::setw(specialization_column_width) << std::left << specialization;
|
|
|
|
// 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)
|
|
std::cout << red;
|
|
|
|
std::cout << line.str() << std::flush;
|
|
}
|
|
|
|
} // namespace progress
|
|
|
|
#ifdef __HIP__
|
|
|
|
/// Kernel that blocks the GPU stream until unblocked or timeout occurs.
|
|
static __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<long long int>(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.
|
|
static __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<double> timeout_duration(timeout_seconds);
|
|
|
|
while(*is_blocked == 1)
|
|
{
|
|
auto elapsed = cuda::std::chrono::duration<double>(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;
|
|
|
|
/// 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_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;
|
|
|
|
// Get device pointers to mapped host memory.
|
|
PRIMBENCH_CHECK(
|
|
host_get_device_pointer(reinterpret_cast<void**>(&temp_device_flag), &m_host_flag, 0));
|
|
PRIMBENCH_CHECK(host_get_device_pointer(reinterpret_cast<void**>(&temp_device_timeout_flag),
|
|
&m_host_timeout_flag,
|
|
0));
|
|
|
|
// 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_CHECK(hipGetDevice(&device_id));
|
|
int wall_clk_rate_k_hz = 0;
|
|
PRIMBENCH_CHECK(
|
|
hipDeviceGetAttribute(&wall_clk_rate_k_hz, hipDeviceAttributeWallClockRate, device_id));
|
|
m_wall_clock_rate = wall_clk_rate_k_hz;
|
|
#endif
|
|
}
|
|
|
|
/// Destructor that unregisters host memory.
|
|
~stream_blocker()
|
|
{
|
|
PRIMBENCH_CHECK(host_unregister(&m_host_flag));
|
|
PRIMBENCH_CHECK(host_unregister(&m_host_timeout_flag));
|
|
}
|
|
|
|
/// Launches a blocking kernel on the stream until unblocked or timed out.
|
|
void block()
|
|
{
|
|
volatile int32_t& flag = m_host_flag;
|
|
flag = 1;
|
|
|
|
volatile int32_t& timeout_flag = m_host_timeout_flag;
|
|
timeout_flag = 0;
|
|
|
|
#ifdef __HIP__
|
|
block_stream_kernel<<<dim3(1), dim3(1), 0, m_stream>>>(m_device_flag,
|
|
m_device_timeout_flag,
|
|
m_stream_blocking_timeout_secs,
|
|
m_wall_clock_rate);
|
|
#elif defined(__CUDACC__)
|
|
block_stream_kernel<<<dim3(1), dim3(1), 0, m_stream>>>(m_device_flag,
|
|
m_device_timeout_flag,
|
|
m_stream_blocking_timeout_secs);
|
|
#endif
|
|
}
|
|
|
|
/// Unblocks the stream by resetting the blocking flag.
|
|
void unblock()
|
|
{
|
|
volatile int32_t& flag = m_host_flag;
|
|
flag = 0;
|
|
}
|
|
|
|
/// Checks if the GPU timed out during blocking.
|
|
///
|
|
/// This should be called after stream synchronization.
|
|
void check_timeout()
|
|
{
|
|
if(m_host_timeout_flag)
|
|
{
|
|
std::cerr
|
|
<< "\nError: Stream blocking timed out after " << m_stream_blocking_timeout_secs
|
|
<< " seconds.\nThe stream is blocked while queueing kernel calls.\nIf your kernel "
|
|
"is synchronous, pass primbench::flags::sync to disable stream blocking.\nYou "
|
|
"can increase the timeout by passing --stream-blocking-timeout-secs <secs>.\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
}
|
|
|
|
private:
|
|
stream_t m_stream;
|
|
double m_stream_blocking_timeout_secs;
|
|
long long int m_wall_clock_rate;
|
|
|
|
int32_t m_host_flag = 0;
|
|
int32_t m_host_timeout_flag = 0;
|
|
|
|
volatile int32_t* m_device_flag = nullptr;
|
|
volatile int32_t* m_device_timeout_flag = nullptr;
|
|
}; // class stream_blocker
|
|
|
|
/// 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;
|
|
using value_type
|
|
= std::variant<std::shared_ptr<json>, std::string, bool, double, int, unsigned int, size_t>;
|
|
|
|
// Using std::map instead of std::unordered_map,
|
|
// as we want a deterministic order.
|
|
using map_type = std::map<key_type, value_type>;
|
|
|
|
/// Adds a key-value pair to the JSON object.
|
|
template<typename T>
|
|
json& add(std::string_view key, T value)
|
|
{
|
|
if constexpr(std::is_same_v<T, json>)
|
|
{
|
|
// Copy nested JSON.
|
|
m_map[std::string(key)] = std::make_shared<json>(value);
|
|
}
|
|
else if constexpr(std::is_convertible_v<T, std::string_view>)
|
|
{
|
|
m_map[std::string(key)] = std::string(std::string_view(value));
|
|
}
|
|
else
|
|
{
|
|
m_map[std::string(key)] = value;
|
|
}
|
|
return *this;
|
|
}
|
|
|
|
/// Serializes the JSON object to a JSON string.
|
|
std::string serialize() const
|
|
{
|
|
std::ostringstream ss;
|
|
ss << "{";
|
|
ss << std::boolalpha;
|
|
|
|
bool is_first = true;
|
|
for(auto const& [key, value] : m_map)
|
|
{
|
|
if(!is_first)
|
|
ss << ",";
|
|
else
|
|
is_first = false;
|
|
|
|
ss << "\"" << key << "\":";
|
|
|
|
std::visit(
|
|
[&](auto const& val)
|
|
{
|
|
using V = std::decay_t<decltype(val)>;
|
|
if constexpr(std::is_same_v<V, std::shared_ptr<json>>)
|
|
{
|
|
ss << val->serialize();
|
|
}
|
|
else if constexpr(std::is_same_v<V, std::string>)
|
|
{
|
|
ss << "\"" << val << "\"";
|
|
}
|
|
else
|
|
{
|
|
ss << val;
|
|
}
|
|
},
|
|
value);
|
|
}
|
|
|
|
ss << "}";
|
|
return ss.str();
|
|
}
|
|
|
|
/// Serializes the JSON object into a human-readable name.
|
|
std::string serialize_name() const
|
|
{
|
|
static const std::vector<std::string_view> blacklist = {"lvl", "algo"};
|
|
static const std::vector<std::string_view> stripped_prefixes = {"rocprim::", "common::"};
|
|
|
|
auto strip_prefixes = [&](std::string& s)
|
|
{
|
|
for(auto p : stripped_prefixes)
|
|
{
|
|
while(true)
|
|
{
|
|
size_t pos = s.find(p);
|
|
if(pos == std::string::npos)
|
|
break;
|
|
s.erase(pos, p.size());
|
|
}
|
|
}
|
|
};
|
|
|
|
// Recursive formatter.
|
|
std::function<std::string(const json&, bool)> fmt
|
|
= [&](const json& obj, bool toplevel) -> std::string
|
|
{
|
|
std::string out;
|
|
if(!toplevel)
|
|
out += "{ ";
|
|
|
|
bool first = true;
|
|
for(auto const& [k_sv, val] : obj.m_map)
|
|
{
|
|
std::string key(k_sv);
|
|
|
|
// Skip blacklisted keys.
|
|
if(std::find(blacklist.begin(), blacklist.end(), key) != blacklist.end())
|
|
continue;
|
|
|
|
std::string vstr;
|
|
std::visit(
|
|
[&](auto const& v)
|
|
{
|
|
using V = std::decay_t<decltype(v)>;
|
|
|
|
if constexpr(std::is_same_v<V, std::shared_ptr<json>>)
|
|
{
|
|
vstr = fmt(*v, false);
|
|
}
|
|
else if constexpr(std::is_same_v<V, std::string>)
|
|
{
|
|
vstr = v;
|
|
strip_prefixes(vstr);
|
|
}
|
|
else
|
|
{
|
|
// Can't use std::to_string(v),
|
|
// as that doesn't trim trailing zeros.
|
|
std::ostringstream ss;
|
|
ss << std::boolalpha;
|
|
ss << v;
|
|
vstr = ss.str();
|
|
}
|
|
},
|
|
val);
|
|
|
|
// Skip default configs.
|
|
if(key == "cfg" && vstr == "default")
|
|
continue;
|
|
|
|
if(!first)
|
|
out += ", ";
|
|
else
|
|
first = false;
|
|
|
|
out += key + ": " + vstr;
|
|
}
|
|
|
|
if(!toplevel)
|
|
out += " }";
|
|
return out;
|
|
};
|
|
|
|
return fmt(*this, true);
|
|
}
|
|
|
|
/// Retrieves a value by key.
|
|
template<typename T>
|
|
const T& get(std::string_view key) const
|
|
{
|
|
const auto& val = m_map.at(std::string(key));
|
|
if constexpr(std::is_same_v<T, json>)
|
|
{
|
|
return *std::get<std::shared_ptr<json>>(val);
|
|
}
|
|
else
|
|
{
|
|
return std::get<T>(val);
|
|
}
|
|
}
|
|
|
|
private:
|
|
map_type m_map = {};
|
|
}; // struct json
|
|
|
|
struct device_storage
|
|
{
|
|
device_storage(size_t size) : m_size(size), m_device_ptr(nullptr)
|
|
{
|
|
if(size > 0)
|
|
PRIMBENCH_CHECK(gpu_malloc(&m_device_ptr, size));
|
|
}
|
|
|
|
~device_storage()
|
|
{
|
|
if(m_device_ptr != nullptr)
|
|
PRIMBENCH_CHECK(gpu_free(m_device_ptr));
|
|
}
|
|
|
|
size_t get_size() const
|
|
{
|
|
return m_size;
|
|
}
|
|
|
|
template<typename T = void>
|
|
T* get_ptr() const
|
|
{
|
|
return static_cast<T*>(m_device_ptr);
|
|
};
|
|
|
|
// This type is not copy assignable.
|
|
device_storage& operator=(const device_storage&) = delete;
|
|
// This type is not copy constructible.
|
|
device_storage(const device_storage&) = delete;
|
|
|
|
private:
|
|
const size_t m_size;
|
|
void* m_device_ptr;
|
|
}; // struct device_storage
|
|
|
|
/// Cache thrashing utility to ensure the cache contains irrelevant data.
|
|
struct cache_thrasher
|
|
{
|
|
public:
|
|
/// 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) {}
|
|
|
|
/// 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_CHECK(
|
|
memset_async(m_device_storage.get_ptr(), 0, m_device_storage.get_size(), stream));
|
|
}
|
|
|
|
// This type is not copy assignable.
|
|
cache_thrasher& operator=(const cache_thrasher&) = delete;
|
|
// This type is not copy constructible.
|
|
cache_thrasher(const cache_thrasher&) = delete;
|
|
|
|
private:
|
|
const device_storage m_device_storage;
|
|
}; // struct cache_thrasher
|
|
|
|
#if PRIMBENCH_HAS_MONITORING
|
|
|
|
/// Warms the GPU, using complex enough dummy kernel code
|
|
/// that the compiler can't optimize it away.
|
|
static __global__
|
|
void warmup_kernel(float* data, int 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)
|
|
{
|
|
x += sinf(x) * cosf(x);
|
|
x *= 1.0000001f;
|
|
x = sqrtf(x + 1.0f);
|
|
x = logf(x + 1.0f);
|
|
}
|
|
|
|
data[idx] = x;
|
|
}
|
|
}
|
|
|
|
class gpu_warmer
|
|
{
|
|
public:
|
|
// Delete all copy/move constructors and assignment operators.
|
|
gpu_warmer(const gpu_warmer&) = delete;
|
|
gpu_warmer& operator=(const gpu_warmer&) = delete;
|
|
gpu_warmer(gpu_warmer&&) = delete;
|
|
gpu_warmer& operator=(gpu_warmer&&) = delete;
|
|
|
|
/// Singleton accessor.
|
|
static gpu_warmer& instance()
|
|
{
|
|
static gpu_warmer instance;
|
|
return instance;
|
|
}
|
|
|
|
gpu_warmer()
|
|
{
|
|
int grid_size;
|
|
PRIMBENCH_CHECK(occupancy_max_potential_block_size(
|
|
&grid_size, // Minimum grid size for full occupancy.
|
|
&m_threads_per_block, // Block size for full occupancy.
|
|
warmup_kernel,
|
|
0, // Dynamic shared memory.
|
|
0 // Block size limit (0 = no limit).
|
|
));
|
|
|
|
m_num_items = grid_size * m_threads_per_block;
|
|
|
|
m_device_storage = std::make_unique<device_storage>(m_num_items * sizeof(float));
|
|
}
|
|
|
|
void warm_up(stream_t stream, uint16_t min_gpu_temp, double max_warming_secs, bool print) const
|
|
{
|
|
auto start = std::chrono::steady_clock::now();
|
|
|
|
while(true)
|
|
{
|
|
uint16_t gpu_temp = monitor::instance().get_temp();
|
|
if(gpu_temp >= min_gpu_temp)
|
|
break;
|
|
|
|
if(use_color() && print)
|
|
progress::print_warming(gpu_temp, min_gpu_temp);
|
|
|
|
dim3 blocks(m_num_items / m_threads_per_block);
|
|
dim3 threads(m_threads_per_block);
|
|
|
|
warmup_kernel<<<blocks, threads, 0, stream>>>(m_device_storage->get_ptr<float>(),
|
|
m_num_items);
|
|
|
|
PRIMBENCH_CHECK(stream_synchronize(stream));
|
|
|
|
auto duration = std::chrono::steady_clock::now() - start;
|
|
if(duration >= std::chrono::duration<double>(max_warming_secs))
|
|
{
|
|
std::cerr << "\nError: Failed to warm up after " << max_warming_secs
|
|
<< " seconds\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
|
|
std::this_thread::sleep_for(std::chrono::milliseconds(1));
|
|
}
|
|
}
|
|
|
|
private:
|
|
int m_threads_per_block;
|
|
int m_num_items;
|
|
|
|
std::unique_ptr<device_storage> m_device_storage;
|
|
}; // class gpu_warmer
|
|
|
|
#endif // PRIMBENCH_HAS_MONITORING
|
|
|
|
/// Manages benchmark execution, GPU warm-up/cool-down, timing, and logging.
|
|
class state
|
|
{
|
|
public:
|
|
state(std::string_view algo,
|
|
json meta,
|
|
size_t specialization_index,
|
|
size_t specialization_count,
|
|
stream_t stream,
|
|
logger& logger,
|
|
stream_blocker& stream_blocker,
|
|
const settings& settings,
|
|
flags::FlagTag flags,
|
|
size_t specialization_column_width,
|
|
size_t index_column_width,
|
|
bool print_index,
|
|
cache_thrasher& cache)
|
|
: stream(stream)
|
|
, size(settings.size)
|
|
, seed(settings.seed)
|
|
, m_algo(algo)
|
|
, m_meta(std::move(meta))
|
|
, m_specialization_index(specialization_index)
|
|
, m_specialization_count(specialization_count)
|
|
, m_logger(logger)
|
|
, m_stream_blocker(stream_blocker)
|
|
, m_settings(settings)
|
|
, m_flags(flags)
|
|
, m_specialization_column_width(specialization_column_width)
|
|
, m_index_column_width(index_column_width)
|
|
, m_print_index(print_index)
|
|
, m_cache(cache)
|
|
{}
|
|
|
|
/// 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)
|
|
{
|
|
std::cerr << "Error: Can't call set_items() twice\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
m_has_set_items = true;
|
|
|
|
m_items = items;
|
|
}
|
|
|
|
/// 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<typename T>
|
|
void add_reads(size_t items)
|
|
{
|
|
if(!m_has_set_items)
|
|
{
|
|
std::cerr << "Error: Can't call add_reads() before calling set_items()\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
if(m_has_set_writes)
|
|
{
|
|
std::cerr << "Error: Can't call add_reads() after calling add_writes()\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
|
|
size_t bytes = items * sizeof(T);
|
|
m_read_write_bytes += bytes;
|
|
}
|
|
|
|
/// 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<typename T>
|
|
void add_writes(size_t items)
|
|
{
|
|
if(!m_has_set_items)
|
|
{
|
|
std::cerr << "Error: Can't call add_writes() before calling set_items()\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
m_has_set_writes = true;
|
|
|
|
size_t bytes = items * sizeof(T);
|
|
m_read_write_bytes += bytes;
|
|
}
|
|
|
|
/// 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<void()> lambda)
|
|
{
|
|
m_run_before_every_iteration_lambda = lambda;
|
|
}
|
|
|
|
/// 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<void()> kernel)
|
|
{
|
|
if(!m_has_set_items)
|
|
{
|
|
std::cerr << "Error: Can't call run() before calling set_items()\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
m_has_run = true;
|
|
|
|
std::string name = m_meta.serialize_name();
|
|
std::string serialized_meta = m_meta.serialize();
|
|
|
|
size_t bytes_per_item = m_read_write_bytes / m_items;
|
|
|
|
#if PRIMBENCH_HAS_MONITORING
|
|
cool_or_warm();
|
|
#endif
|
|
|
|
init_kernels_per_batch(kernel);
|
|
|
|
// Reserve space for start and stop events for each iteration.
|
|
std::vector<event_t> events(m_kernels_per_batch * 2);
|
|
for(auto& event : events)
|
|
PRIMBENCH_CHECK(event_create(&event));
|
|
|
|
uint64_t iterations = 0;
|
|
std::vector<float> iterations_ms(m_kernels_per_batch);
|
|
|
|
double elapsed_gpu_secs = 0.0;
|
|
|
|
double estimated_remaining_secs = get_estimated_remaining_secs();
|
|
|
|
#if PRIMBENCH_HAS_MONITORING
|
|
monitor::instance().save_start_temp();
|
|
#endif
|
|
|
|
auto start = std::chrono::steady_clock::now();
|
|
|
|
while(true)
|
|
{
|
|
if(run_iteration(kernel,
|
|
events,
|
|
iterations,
|
|
iterations_ms,
|
|
start,
|
|
elapsed_gpu_secs,
|
|
name,
|
|
serialized_meta,
|
|
bytes_per_item,
|
|
estimated_remaining_secs))
|
|
{
|
|
break;
|
|
}
|
|
}
|
|
|
|
for(const auto& event : events)
|
|
PRIMBENCH_CHECK(event_destroy(event));
|
|
}
|
|
|
|
/// Registers \ref test_lambda, which runs once during warmup.
|
|
///
|
|
/// Must not be called after \ref run().
|
|
///
|
|
/// Define `PRIMBENCH_NO_TEST` to disable.
|
|
void test(std::function<void()> test_lambda)
|
|
{
|
|
if(m_has_run)
|
|
{
|
|
std::cerr << "Error: Can't call test() after calling run()\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
|
|
#ifndef PRIMBENCH_NO_TEST
|
|
m_test_lambda = test_lambda;
|
|
#endif
|
|
}
|
|
|
|
// 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.
|
|
|
|
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<void()>& kernel,
|
|
std::vector<event_t>& events,
|
|
uint64_t& iterations,
|
|
std::vector<float>& iterations_ms,
|
|
const std::chrono::steady_clock::time_point& start,
|
|
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;
|
|
|
|
#if PRIMBENCH_HAS_MONITORING
|
|
cool_or_warm(false);
|
|
#endif
|
|
|
|
iterations++;
|
|
|
|
run_batch(events, kernel);
|
|
|
|
fill_iterations_ms(iterations_ms, events);
|
|
|
|
double batch_gpu_ms = std::accumulate(iterations_ms.begin(), iterations_ms.end(), 0.0);
|
|
m_times.emplace_back(batch_gpu_ms);
|
|
|
|
m_logger.save(batch_gpu_ms, iterations_ms);
|
|
|
|
// Gather batch_window_size number of iteration times.
|
|
auto window_start = m_times.end() - std::min(iterations, s.batch_window_size);
|
|
std::vector<double> recent_times(window_start, m_times.end());
|
|
|
|
// 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;
|
|
|
|
double batch_gpu_secs = batch_gpu_ms / 1000;
|
|
|
|
elapsed_gpu_secs += batch_gpu_secs;
|
|
|
|
double bytes_per_batch = m_read_write_bytes * m_kernels_per_batch;
|
|
double bytes_per_sec = bytes_per_batch / batch_gpu_secs;
|
|
|
|
double items_per_batch = m_items * m_kernels_per_batch;
|
|
double items_per_sec = items_per_batch / batch_gpu_secs;
|
|
|
|
auto now = std::chrono::steady_clock::now();
|
|
auto elapsed_host = now - start;
|
|
|
|
// Stop early if the noise stabilized.
|
|
bool stop_early = elapsed_host >= std::chrono::duration<double>(s.min_secs)
|
|
&& iterations >= s.batch_window_size
|
|
&& noise_percent < s.noise_tolerance_percent;
|
|
|
|
// Stop early if the benchmark has been noisy for too long.
|
|
bool noise_timeout = elapsed_host > std::chrono::duration<double>(s.noise_timeout_secs)
|
|
&& iterations >= s.batch_window_size
|
|
&& noise_percent >= s.noise_tolerance_percent;
|
|
|
|
double elapsed_host_secs = std::chrono::duration<double>(elapsed_host).count();
|
|
|
|
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";
|
|
|
|
// Only print progress if it's a terminal (real-time updates)
|
|
// or if the specialization has finished (single final update).
|
|
if(use_color() || stop_early || noise_timeout)
|
|
{
|
|
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,
|
|
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,
|
|
elapsed_host_secs,
|
|
elapsed_gpu_secs,
|
|
noise_timeout);
|
|
|
|
return true;
|
|
}
|
|
|
|
return false;
|
|
}
|
|
|
|
#if PRIMBENCH_HAS_MONITORING
|
|
/// Gets the GPU within a stable temperature range.
|
|
void cool_or_warm(bool print = true) const
|
|
{
|
|
const auto& s = m_settings;
|
|
|
|
// Warm up GPU.
|
|
gpu_warmer::instance().warm_up(stream, s.min_gpu_temp, s.max_warming_secs, print);
|
|
|
|
// Wait for GPU to cool down.
|
|
auto cool_down_start_time = std::chrono::steady_clock::now();
|
|
while(true)
|
|
{
|
|
uint16_t gpu_temp = monitor::instance().get_temp();
|
|
if(gpu_temp <= s.max_gpu_temp)
|
|
break;
|
|
|
|
if(use_color() && print)
|
|
progress::print_cooling(gpu_temp, s.max_gpu_temp);
|
|
|
|
auto duration = std::chrono::steady_clock::now() - cool_down_start_time;
|
|
if(duration >= std::chrono::duration<double>(s.max_cooling_secs))
|
|
{
|
|
std::cerr << "\nError: Failed to cool down after " << s.max_cooling_secs
|
|
<< " seconds\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
|
|
std::this_thread::sleep_for(std::chrono::milliseconds(100));
|
|
}
|
|
}
|
|
#endif
|
|
|
|
/// Determines number of kernels per batch based on minimum duration.
|
|
void init_kernels_per_batch(std::function<void()> kernel)
|
|
{
|
|
std::vector<event_t> events(2);
|
|
std::vector<float> iterations_ms;
|
|
m_kernels_per_batch = 1;
|
|
|
|
// Without this, the very first timed batch can be very slow.
|
|
log("Running warmup");
|
|
for(auto& event : events)
|
|
PRIMBENCH_CHECK(event_create(&event));
|
|
run_batch(events, kernel);
|
|
if(m_test_lambda)
|
|
{
|
|
primbench::log("Running tests");
|
|
m_test_lambda();
|
|
}
|
|
for(const auto& event : events)
|
|
PRIMBENCH_CHECK(event_destroy(event));
|
|
|
|
while(true)
|
|
{
|
|
log("Timing batch size ", m_kernels_per_batch);
|
|
|
|
// Reserve space for start and stop events for each iteration.
|
|
events.resize(m_kernels_per_batch * 2);
|
|
|
|
iterations_ms.resize(m_kernels_per_batch);
|
|
|
|
for(auto& event : events)
|
|
PRIMBENCH_CHECK(event_create(&event));
|
|
|
|
run_batch(events, kernel);
|
|
|
|
fill_iterations_ms(iterations_ms, events);
|
|
|
|
m_ms_per_batch = std::accumulate(iterations_ms.begin(), iterations_ms.end(), 0.0);
|
|
std::chrono::duration<double> batch_ms(m_ms_per_batch);
|
|
|
|
for(const auto& event : events)
|
|
PRIMBENCH_CHECK(event_destroy(event));
|
|
|
|
if(batch_ms > std::chrono::duration<double>(m_settings.min_gpu_ms_per_batch))
|
|
break;
|
|
|
|
// Doubling is the simplest form of exponential growth.
|
|
m_kernels_per_batch *= 2;
|
|
}
|
|
}
|
|
|
|
/// Executes a batch of kernel iterations with event timing.
|
|
void run_batch(const std::vector<event_t>& events, std::function<void()> kernel)
|
|
{
|
|
for(size_t i = 0; i < m_kernels_per_batch; i++)
|
|
{
|
|
if(m_run_before_every_iteration_lambda)
|
|
m_run_before_every_iteration_lambda();
|
|
|
|
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, 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_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_CHECK(event_record(events[i * 2 + 1], stream));
|
|
|
|
// Allows the GPU to start running its queued events.
|
|
if(!m_flags.has(flags::Flags::sync))
|
|
m_stream_blocker.unblock();
|
|
|
|
// Catches kernel launch errors.
|
|
// 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 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.
|
|
// 64 was chosen empirically. It'll sync on i=63, i=127, etc.
|
|
constexpr size_t n = 64;
|
|
if(i % n == n - 1)
|
|
{
|
|
PRIMBENCH_CHECK(stream_synchronize(stream));
|
|
// Catch runtime/device execution errors.
|
|
PRIMBENCH_CHECK(get_last_error());
|
|
|
|
// Check for blocking kernel timeout after sync.
|
|
if(!m_flags.has(flags::Flags::sync))
|
|
m_stream_blocker.check_timeout();
|
|
}
|
|
}
|
|
|
|
// Final stream synchronization.
|
|
PRIMBENCH_CHECK(stream_synchronize(stream));
|
|
// Catch any runtime/device errors from the last kernels.
|
|
PRIMBENCH_CHECK(get_last_error());
|
|
|
|
// Final blocking kernel timeout check.
|
|
if(!m_flags.has(flags::Flags::sync))
|
|
m_stream_blocker.check_timeout();
|
|
}
|
|
|
|
/// Fills iteration times (ms) using HIP event timing.
|
|
void fill_iterations_ms(std::vector<float>& iterations_ms,
|
|
const std::vector<event_t>& 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_CHECK(event_elapsed_time(&iteration_ms, events[i * 2], events[i * 2 + 1]));
|
|
|
|
iterations_ms[i] = iteration_ms;
|
|
}
|
|
}
|
|
|
|
/// Clears GPU caches.
|
|
void clear_gpu_cache(stream_t stream) const
|
|
{
|
|
m_cache.clear_cache(stream);
|
|
}
|
|
|
|
/// 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<double>(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<double>& times) const
|
|
{
|
|
return std::reduce(times.begin(), times.end()) / times.size();
|
|
}
|
|
|
|
/// Computes standard deviation of time samples.
|
|
double get_stddev(const std::vector<double>& times) const
|
|
{
|
|
const size_t n = times.size();
|
|
if(n <= 1)
|
|
return 0.0;
|
|
|
|
double mean = std::accumulate(times.begin(), times.end(), 0.0) / n;
|
|
double sum_sq = 0.0;
|
|
for(double x : times)
|
|
sum_sq += (x - mean) * (x - mean);
|
|
|
|
return std::sqrt(sum_sq / (n - 1));
|
|
}
|
|
|
|
/// Computes coefficient of variation (CV).
|
|
double get_cv(const std::vector<double>& times, double stddev, double mean) const
|
|
{
|
|
return times.size() >= 2 ? stddev / mean : 0.0;
|
|
}
|
|
|
|
std::string m_algo;
|
|
const json m_meta;
|
|
|
|
size_t m_specialization_index;
|
|
size_t m_specialization_count;
|
|
|
|
logger& m_logger;
|
|
stream_blocker& m_stream_blocker;
|
|
|
|
const settings& m_settings;
|
|
|
|
flags::FlagTag m_flags;
|
|
|
|
size_t m_specialization_column_width;
|
|
|
|
size_t m_index_column_width;
|
|
bool m_print_index;
|
|
|
|
cache_thrasher& m_cache;
|
|
|
|
std::function<void()> m_run_before_every_iteration_lambda = nullptr;
|
|
std::function<void()> m_test_lambda = nullptr;
|
|
std::vector<double> m_times;
|
|
size_t m_kernels_per_batch = 0;
|
|
double m_ms_per_batch = 0.0;
|
|
|
|
bool m_has_set_items = false;
|
|
bool m_has_set_writes = false;
|
|
bool m_has_run = false;
|
|
size_t m_items = 0;
|
|
size_t m_read_write_bytes = 0;
|
|
}; // class state
|
|
|
|
/// Simple command-line argument parser.
|
|
class cli
|
|
{
|
|
public:
|
|
cli(int argc, char** argv) : _appname(argv[0])
|
|
{
|
|
// Always register --help.
|
|
register_description("help", "Display this help and exit.");
|
|
|
|
for(int i = 1; i < argc; ++i)
|
|
{
|
|
std::string arg = argv[i];
|
|
|
|
if(arg.rfind("--", 0) != 0)
|
|
continue;
|
|
|
|
std::string key = arg.substr(2);
|
|
std::string value = "";
|
|
|
|
// Consume all consecutive arguments that don't start with `--`.
|
|
while(i + 1 < argc && std::string(argv[i + 1]).rfind("--", 0) != 0)
|
|
{
|
|
if(!value.empty())
|
|
value += " ";
|
|
value += argv[++i];
|
|
}
|
|
|
|
_parsed[key] = value;
|
|
}
|
|
}
|
|
|
|
/// Gets argument value, registering it with default and description if needed.
|
|
template<typename T>
|
|
T get(std::string_view name, const T& default_val, std::string_view description)
|
|
{
|
|
std::string key{name};
|
|
|
|
// 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".
|
|
if constexpr(std::is_same_v<T, bool>)
|
|
{
|
|
oss << std::boolalpha << default_val;
|
|
}
|
|
// For vectors, join elements with spaces.
|
|
else if constexpr(is_vector<T>::value)
|
|
{
|
|
for(size_t i = 0; i < default_val.size(); ++i)
|
|
{
|
|
if(i > 0)
|
|
oss << " ";
|
|
oss << default_val[i];
|
|
}
|
|
}
|
|
else
|
|
{
|
|
oss << default_val;
|
|
}
|
|
|
|
_defaults[key] = oss.str();
|
|
_registered.insert(key);
|
|
|
|
// For bools, ensure no value was provided on command line.
|
|
if constexpr(std::is_same_v<T, bool>)
|
|
{
|
|
auto it = _parsed.find(key);
|
|
if(it != _parsed.end() && !it->second.empty())
|
|
{
|
|
std::cerr << "Error: Boolean flag --" << key << " does not take a value.\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
}
|
|
}
|
|
|
|
auto it = _parsed.find(key);
|
|
auto default_it = _defaults.find(key);
|
|
|
|
// If not provided on command line, use default.
|
|
if(it == _parsed.end())
|
|
{
|
|
T out{};
|
|
if(default_it != _defaults.end() && !default_it->second.empty())
|
|
{
|
|
std::istringstream ss(default_it->second);
|
|
|
|
// For bools, use boolalpha to parse "true"/"false".
|
|
if constexpr(std::is_same_v<T, bool>)
|
|
{
|
|
ss >> std::boolalpha >> out;
|
|
if(ss.fail())
|
|
{
|
|
std::cerr << "Error: Failed to parse default for --" << key
|
|
<< ": invalid value \"" << default_it->second << "\"\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
}
|
|
else if constexpr(is_vector<T>::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
|
|
{
|
|
ss >> out;
|
|
if(!ss || !ss.eof())
|
|
{
|
|
std::cerr << "Error: Failed to parse default for --" << key
|
|
<< ": invalid value \"" << default_it->second << "\"\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
}
|
|
}
|
|
return out;
|
|
}
|
|
|
|
// Parse provided value.
|
|
if constexpr(std::is_same_v<T, bool>)
|
|
{
|
|
return it->second.empty() ? true : false;
|
|
}
|
|
else if constexpr(is_vector<T>::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<Elem, std::string>)
|
|
{
|
|
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{};
|
|
std::istringstream ss(it->second);
|
|
|
|
if constexpr(std::is_same_v<T, std::string>)
|
|
{
|
|
std::getline(ss, out);
|
|
}
|
|
else
|
|
{
|
|
ss >> out;
|
|
}
|
|
|
|
if(!ss)
|
|
{
|
|
std::cerr << "Error: Failed to parse --" << key << ": invalid value \""
|
|
<< it->second << "\"\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
|
|
// When "--size" is specified, the number can optionally be suffixed with KiB/MiB/GiB.
|
|
if(key == "size")
|
|
{
|
|
if constexpr(std::is_same_v<T, size_t>)
|
|
{
|
|
std::string suffix;
|
|
ss >> suffix;
|
|
|
|
if(!suffix.empty())
|
|
{
|
|
if(suffix == "KiB")
|
|
out *= KiB;
|
|
else if(suffix == "MiB")
|
|
out *= MiB;
|
|
else if(suffix == "GiB")
|
|
out *= GiB;
|
|
else
|
|
{
|
|
std::cerr << "Error: Failed to parse --" << key << ": unknown suffix \""
|
|
<< suffix << "\"\n";
|
|
std::exit(EXIT_FAILURE);
|
|
}
|
|
}
|
|
}
|
|
}
|
|
|
|
return out;
|
|
}
|
|
}
|
|
|
|
/// Returns all registered arguments with their parsed values.
|
|
std::map<std::string, settings::custom_arg_value> get_all_custom_options() const
|
|
{
|
|
std::map<std::string, settings::custom_arg_value> custom_args;
|
|
|
|
// Skip built-in arguments that are already in settings.
|
|
static const std::unordered_set<std::string> builtin_args
|
|
= {"help",
|
|
"size",
|
|
"hot",
|
|
"seed",
|
|
"json-out",
|
|
"csv-out",
|
|
"filter",
|
|
"dry",
|
|
"min-gpu-ms-per-batch",
|
|
"min-secs",
|
|
"noise-timeout-secs",
|
|
"batch-window-size",
|
|
"noise-tolerance-percent",
|
|
"min-gpu-temp",
|
|
"max-gpu-temp",
|
|
"max-warming-secs",
|
|
"max-cooling-secs",
|
|
"output-batches",
|
|
"spaces-per-indent",
|
|
"stream-blocking-timeout-secs"};
|
|
|
|
auto parse_value = [](const std::string& value) -> settings::custom_arg_value
|
|
{
|
|
if(value.empty())
|
|
{
|
|
return true; // Boolean flag.
|
|
}
|
|
|
|
// Check for boolean strings.
|
|
if(value == "true")
|
|
{
|
|
return true;
|
|
}
|
|
if(value == "false")
|
|
{
|
|
return false;
|
|
}
|
|
|
|
// Try int.
|
|
try
|
|
{
|
|
size_t idx = 0;
|
|
int int_val = std::stoi(value, &idx);
|
|
if(idx == value.size())
|
|
{
|
|
return int_val;
|
|
}
|
|
}
|
|
catch(...)
|
|
{}
|
|
|
|
// Try double.
|
|
try
|
|
{
|
|
size_t idx = 0;
|
|
double double_val = std::stod(value, &idx);
|
|
if(idx == value.size())
|
|
{
|
|
return double_val;
|
|
}
|
|
}
|
|
catch(...)
|
|
{}
|
|
|
|
// Fall back to string.
|
|
return value;
|
|
};
|
|
|
|
// Process all registered arguments, checking both defaults and parsed values.
|
|
std::unordered_set<std::string> processed;
|
|
|
|
// First, add all parsed custom arguments.
|
|
for(const auto& [key, value] : _parsed)
|
|
{
|
|
if(builtin_args.find(key) == builtin_args.end())
|
|
{
|
|
custom_args[key] = parse_value(value);
|
|
processed.insert(key);
|
|
}
|
|
}
|
|
|
|
// 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()
|
|
&& processed.find(key) == processed.end())
|
|
{
|
|
custom_args[key] = parse_value(default_val);
|
|
}
|
|
}
|
|
|
|
return custom_args;
|
|
}
|
|
|
|
/// Prints help if requested and validates all arguments are registered.
|
|
void finalize() const
|
|
{
|
|
possibly_print_help();
|
|
validate_arguments();
|
|
}
|
|
|
|
private:
|
|
/// Prints help message and exits if --help was requested.
|
|
void possibly_print_help() const
|
|
{
|
|
if(_parsed.find("help") == _parsed.end())
|
|
return;
|
|
|
|
std::cout << "Usage: " << _appname << " [options]\n\n";
|
|
for(const auto& [key, desc] : _descriptions)
|
|
{
|
|
std::cout << " --" << key;
|
|
|
|
// Print default if available.
|
|
auto it_def = _defaults.find(key);
|
|
if(it_def != _defaults.end() && !it_def->second.empty())
|
|
{
|
|
// Boolean flags are implicitly false by default.
|
|
if(it_def->second != "false")
|
|
{
|
|
std::cout << " (default: " << it_def->second << ")";
|
|
}
|
|
}
|
|
std::cout << "\n";
|
|
|
|
if(!desc.empty())
|
|
{
|
|
std::cout << " " << desc << "\n";
|
|
}
|
|
}
|
|
exit(EXIT_SUCCESS);
|
|
}
|
|
|
|
/// Validates that all parsed arguments were registered; exits if not.
|
|
void validate_arguments() const
|
|
{
|
|
for(const auto& [key, value] : _parsed)
|
|
{
|
|
if(_registered.find(key) == _registered.end())
|
|
{
|
|
std::cerr << "Error: Unrecognized argument --" << key << ".\n";
|
|
std::cerr
|
|
<< "To register this argument, call .get() with a type and description:\n";
|
|
std::cerr << " executor.get<std::string>(\"" << key
|
|
<< "\", default_value, \"Description\");\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
}
|
|
}
|
|
|
|
std::string _appname; ///< Stores argv[0].
|
|
std::unordered_map<std::string, std::string>
|
|
_parsed; ///< Stores pairs of passed flag_name+flag_value.
|
|
|
|
std::vector<std::pair<std::string, std::string>> _descriptions; ///< Preserves insertion order.
|
|
|
|
std::unordered_set<std::string>
|
|
_description_keys_set; ///< Prevents duplicate descriptions being printed.
|
|
|
|
std::unordered_map<std::string, std::string>
|
|
_defaults; ///< Stores string representation of default values.
|
|
|
|
std::unordered_set<std::string> _registered; ///< Tracks which arguments were registered.
|
|
|
|
/// 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};
|
|
|
|
// All descriptions should be sentences ending with a period.
|
|
if(!desc.empty() && desc.back() != '.')
|
|
{
|
|
desc.push_back('.');
|
|
}
|
|
|
|
_descriptions.emplace_back(key, desc);
|
|
_description_keys_set.insert(key);
|
|
}
|
|
|
|
// Helper to detect vectors.
|
|
template<typename T>
|
|
struct is_vector : std::false_type
|
|
{};
|
|
|
|
template<typename T, typename A>
|
|
struct is_vector<std::vector<T, A>> : std::true_type
|
|
{};
|
|
}; // class cli
|
|
|
|
} // namespace detail
|
|
|
|
using json = detail::json;
|
|
using state = detail::state;
|
|
|
|
/// Used to retrieve the name of a type that was registered with PRIMBENCH_REGISTER_TYPE().
|
|
template<class T>
|
|
std::string name()
|
|
{
|
|
return detail::type_name<T>::name;
|
|
}
|
|
|
|
/// 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<typename... Args>
|
|
void log(Args&&... args)
|
|
{
|
|
if(!detail::use_color())
|
|
return;
|
|
|
|
std::cout << detail::clearline << detail::gray;
|
|
(std::cout << ... << args);
|
|
std::cout << detail::reset << std::flush;
|
|
}
|
|
|
|
/// Namespace for flag definitions and utilities.
|
|
namespace flags
|
|
{
|
|
|
|
/// FlagTag representing no flags.
|
|
inline constexpr detail::flags::FlagTag none{detail::flags::Flags::none};
|
|
|
|
/// FlagTag representing the sync flag.
|
|
inline constexpr detail::flags::FlagTag sync{detail::flags::Flags::sync};
|
|
|
|
} // namespace flags
|
|
|
|
/// 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
|
|
{
|
|
/// 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;
|
|
|
|
/// 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;
|
|
};
|
|
|
|
/// Generates and manages `SeedCount` seeds, given a single input seed.
|
|
template<size_t SeedCount>
|
|
class seeds
|
|
{
|
|
public:
|
|
seeds() = delete;
|
|
|
|
seeds(uint32_t seed)
|
|
{
|
|
std::seed_seq seq{seed};
|
|
seq.generate(m_seeds.begin(), m_seeds.end());
|
|
}
|
|
|
|
uint32_t operator[](size_t index) const
|
|
{
|
|
return m_seeds[index];
|
|
}
|
|
|
|
private:
|
|
std::array<uint32_t, SeedCount> m_seeds;
|
|
}; // class seeds
|
|
|
|
/// 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:
|
|
/// Constructs the executor, and runs setup code.
|
|
executor(int argc,
|
|
char* argv[],
|
|
primbench::settings settings = {},
|
|
detail::flags::FlagTag flags = flags::none,
|
|
stream_t stream = default_stream)
|
|
: m_settings(settings)
|
|
, m_flags(flags)
|
|
, m_stream(stream)
|
|
, m_own_stream(stream == default_stream)
|
|
, m_cli(argc, argv)
|
|
{
|
|
get_logger().save_program_start_time();
|
|
|
|
parse();
|
|
|
|
// If user did not provide a stream, create a fast private one.
|
|
// We can't use default_stream, as it synchronizes with the host.
|
|
if(m_own_stream)
|
|
PRIMBENCH_CHECK(detail::stream_create(&m_stream));
|
|
|
|
m_stream_blocker
|
|
= std::make_unique<detail::stream_blocker>(m_stream,
|
|
m_settings.stream_blocking_timeout_secs);
|
|
}
|
|
|
|
~executor()
|
|
{
|
|
if(m_own_stream && m_stream != nullptr)
|
|
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, which allows the function to be called in global scope.
|
|
template<typename Benchmark, typename... Args>
|
|
static bool queue(Args&&... args)
|
|
{
|
|
specializations.push_back(std::make_unique<Benchmark>(std::forward<Args>(args)...));
|
|
return true;
|
|
}
|
|
|
|
/// \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<typename BulkCreateFunction>
|
|
static bool queue_autotune(BulkCreateFunction&& fn)
|
|
{
|
|
std::forward<BulkCreateFunction>(fn)(specializations);
|
|
return true;
|
|
}
|
|
|
|
/// Prepares and runs all queued benchmark specializations.
|
|
void run()
|
|
{
|
|
static bool run_called = false;
|
|
if(run_called)
|
|
{
|
|
std::cerr << "Error: executor's run() can't be called more than once per algorithm\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
run_called = true;
|
|
|
|
m_cli.finalize();
|
|
|
|
// Capture custom arguments after all have been registered.
|
|
m_settings.custom_args = m_cli.get_all_custom_options();
|
|
|
|
filter_specializations();
|
|
ensure_specializations_exist();
|
|
sort_specializations();
|
|
|
|
std::string algorithm = get_common_algorithm();
|
|
|
|
get_logger().init(algorithm, specializations.size(), m_settings, m_flags);
|
|
|
|
m_specialization_column_width = compute_max_specialization_width(algorithm);
|
|
|
|
constexpr size_t min_width = sizeof("9h 59m 59s") - 1;
|
|
|
|
m_index_column_width = std::max(min_width,
|
|
std::string("Index/").size()
|
|
+ std::to_string(specializations.size()).size());
|
|
|
|
m_print_index = specializations.size() > 1;
|
|
|
|
print_header(algorithm);
|
|
|
|
run_all_specializations();
|
|
|
|
get_logger().output_summary();
|
|
}
|
|
|
|
/// Parses a command-line argument.
|
|
template<typename T>
|
|
T get(std::string_view name, const T& default_val, std::string_view description)
|
|
{
|
|
return m_cli.get<T>(name, default_val, description);
|
|
}
|
|
|
|
private:
|
|
/// Parse optional arguments.
|
|
void parse()
|
|
{
|
|
auto& cli = m_cli;
|
|
auto& s = m_settings;
|
|
|
|
s.size = cli.get<size_t>("size",
|
|
s.size,
|
|
"Input size. Benchmarks decide what this represents, but it is "
|
|
"commonly the number of bytes or items."
|
|
" Supports the suffixes KiB/MiB/GiB, e.g. `--size 256KiB`.");
|
|
if(s.size == 0)
|
|
{
|
|
std::cerr << "Error: --size must be greater than 0\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
|
|
s.hot
|
|
= cli.get<bool>("hot", s.hot, "Skip clearing the GPU cache between batch iterations.");
|
|
|
|
s.seed = cli.get<uint32_t>("seed", s.seed, "Seed used for input generation.");
|
|
|
|
s.json_out = cli.get<std::string>("json-out",
|
|
s.json_out,
|
|
"JSON path to write benchmark results to.");
|
|
|
|
s.csv_out
|
|
= cli.get<std::string>("csv-out", s.csv_out, "CSV path to write benchmark results to.");
|
|
|
|
s.filter = cli.get<std::string>("filter",
|
|
s.filter,
|
|
"Regex filter of specialization names to benchmark.");
|
|
|
|
s.dry = cli.get<bool>("dry",
|
|
s.dry,
|
|
"Perform a dry run. JSON and CSV files are still output, but "
|
|
"benchmark setup and execution are skipped.");
|
|
|
|
s.min_gpu_ms_per_batch
|
|
= cli.get<double>("min-gpu-ms-per-batch",
|
|
s.min_gpu_ms_per_batch,
|
|
"Minimum duration of a batch in milliseconds (GPU time).");
|
|
if(s.min_gpu_ms_per_batch <= 0.0)
|
|
{
|
|
std::cerr << "Error: --min-gpu-ms-per-batch must be greater than 0\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
|
|
s.min_secs = cli.get<double>("min-secs",
|
|
s.min_secs,
|
|
"Minimum total benchmark duration in seconds (wall time).");
|
|
if(s.min_secs <= 0.0)
|
|
{
|
|
std::cerr << "Error: --min-secs must be greater than 0\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
|
|
s.noise_timeout_secs
|
|
= cli.get<double>("noise-timeout-secs",
|
|
s.noise_timeout_secs,
|
|
"Maximum total benchmark duration in seconds before timing out a "
|
|
"noisy run (wall time).");
|
|
if(s.noise_timeout_secs <= 0.0)
|
|
{
|
|
std::cerr << "Error: --noise-timeout-secs must be greater than 0\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
if(s.min_secs > s.noise_timeout_secs)
|
|
{
|
|
std::cerr << "Error: --min-secs must be equal to or less than --noise-timeout-secs\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
|
|
s.batch_window_size
|
|
= cli.get<size_t>("batch-window-size",
|
|
s.batch_window_size,
|
|
"Number of batch times used in the noise (coefficient of variation) "
|
|
"window to decide early benchmark stopping.");
|
|
if(s.batch_window_size == 0)
|
|
{
|
|
std::cerr << "Error: --batch-window-size must be greater than 0\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
|
|
s.noise_tolerance_percent
|
|
= cli.get<double>("noise-tolerance-percent",
|
|
s.noise_tolerance_percent,
|
|
"Noise tolerance of batch times in percent, used to determine "
|
|
"whether a benchmark can be stopped early.");
|
|
if(s.noise_tolerance_percent <= 0.0)
|
|
{
|
|
std::cerr << "Error: --noise-tolerance-percent must be greater than 0\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
|
|
s.min_gpu_temp = cli.get<uint16_t>(
|
|
"min-gpu-temp",
|
|
s.min_gpu_temp,
|
|
"Minimum GPU temperature in °C. Too low slows benchmarks; too high increases noise.");
|
|
s.max_gpu_temp = cli.get<uint16_t>(
|
|
"max-gpu-temp",
|
|
s.max_gpu_temp,
|
|
"Maximum GPU temperature in °C. Too low slows benchmarks; too high increases noise.");
|
|
if(s.min_gpu_temp > s.max_gpu_temp)
|
|
{
|
|
std::cerr << "Error: --min-gpu-temp must be equal to or less than --max-gpu-temp\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
|
|
s.max_warming_secs
|
|
= cli.get<double>("max-warming-secs",
|
|
s.max_warming_secs,
|
|
"Maximum seconds allowed for GPU warming before an error is thrown.");
|
|
if(s.max_warming_secs <= 0.0)
|
|
{
|
|
std::cerr << "Error: --max-warming-secs must be greater than 0\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
|
|
s.max_cooling_secs
|
|
= cli.get<double>("max-cooling-secs",
|
|
s.max_cooling_secs,
|
|
"Maximum seconds allowed for GPU cooling before an error is thrown.");
|
|
if(s.max_cooling_secs <= 0.0)
|
|
{
|
|
std::cerr << "Error: --max-cooling-secs must be greater than 0\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
|
|
s.output_batches = cli.get<bool>(
|
|
"output-batches",
|
|
s.output_batches,
|
|
"Output a `batches` array for each specialization, containing per-batch details.");
|
|
|
|
s.spaces_per_indent = cli.get<uint32_t>(
|
|
"spaces-per-indent",
|
|
s.spaces_per_indent,
|
|
"Number of spaces per indentation level in JSON output. Set to 0 for no indentation.");
|
|
if(s.spaces_per_indent > 8)
|
|
{
|
|
std::cerr << "Error: --spaces-per-indent must be less than or equal to 8\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
|
|
s.stream_blocking_timeout_secs = cli.get<double>(
|
|
"stream-blocking-timeout-secs",
|
|
s.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.");
|
|
if(s.stream_blocking_timeout_secs <= 0.0)
|
|
{
|
|
std::cerr << "Error: --stream-blocking-timeout-secs must be greater than 0\n";
|
|
exit(EXIT_FAILURE);
|
|
}
|
|
}
|
|
|
|
/// 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<std::string>("algo");
|
|
|
|
for(const auto& bp : specializations)
|
|
{
|
|
auto bp_algo = bp->meta().get<std::string>("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<std::string> 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<std::string>("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),
|
|
specialization_index,
|
|
specializations.size(),
|
|
m_stream,
|
|
get_logger(),
|
|
*m_stream_blocker,
|
|
m_settings,
|
|
m_flags,
|
|
m_specialization_column_width,
|
|
m_index_column_width,
|
|
m_print_index,
|
|
m_cache);
|
|
}
|
|
|
|
/// 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();
|
|
|
|
// A dry run doesn't run anything on the GPU, so output placeholder values.
|
|
size_t kernels_per_batch = 0;
|
|
double ms_per_batch = 0.0;
|
|
double bytes_per_sec = 0.0;
|
|
double items_per_sec = 0.0;
|
|
double bytes_per_item = 0.0;
|
|
size_t items = 0;
|
|
double noise_percent = 0.0;
|
|
double elapsed_host_secs = 0.0;
|
|
double elapsed_gpu_secs = 0.0;
|
|
bool noise_timeout = false;
|
|
|
|
detail::progress::print_dry_progress(name,
|
|
algo,
|
|
specialization_index,
|
|
m_specialization_column_width,
|
|
m_index_column_width,
|
|
m_print_index);
|
|
|
|
get_logger().output_specialization(specialization_index,
|
|
name,
|
|
serialized_meta,
|
|
kernels_per_batch,
|
|
ms_per_batch,
|
|
bytes_per_sec,
|
|
items_per_sec,
|
|
bytes_per_item,
|
|
items,
|
|
noise_percent,
|
|
elapsed_host_secs,
|
|
elapsed_gpu_secs,
|
|
noise_timeout);
|
|
}
|
|
|
|
/// Returns the logger singleton.
|
|
detail::logger& get_logger()
|
|
{
|
|
return detail::logger::instance();
|
|
}
|
|
|
|
using specializations_t = std::vector<std::unique_ptr<benchmark_interface>>;
|
|
|
|
/// This vector is static, as benchmarks can be registered in global scope.
|
|
inline static specializations_t specializations;
|
|
|
|
settings m_settings; ///< CLI user settings.
|
|
|
|
detail::flags::FlagTag m_flags; ///< Executor flags.
|
|
|
|
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.
|
|
|
|
std::unique_ptr<detail::stream_blocker>
|
|
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.
|
|
|
|
detail::cache_thrasher m_cache; ///< Cache clearing utility.
|
|
}; // class executor
|
|
|
|
} // namespace primbench
|