-
Notifications
You must be signed in to change notification settings - Fork 11
Contributing: Writing Benchmarks
Every operator should ship with a benchmark so its performance is tracked under the same harness as the rest of the suite. This page covers how to write a benchmark unit — file layout, the macros, parameter passing, and the helpers you should reuse. For methodology (what the suite measures, how it filters noise, how to interpret results) and analysis tooling, see benchmarks/README.md.
One file per operator at benchmarks/src/roccv/bench_<name>.cpp. The benchmark CMake target uses GLOB over benchmarks/src/roccv/*.cpp, so dropping a new file into that directory is enough — re-run CMake configure and it will be picked up. There are no CMakeLists.txt edits.
The shared helpers live in benchmarks/src/roccv/roccv_bench_helpers.hpp, and the framework headers (registry, parameter passing, timing) live under benchmarks/roccvbench/include/roccvbench/.
Every rocCV benchmark file follows the same three-part shape:
- A templated runner function that builds tensors, fills them with deterministic random data, and calls
RecordRunsaround the operator invocation. - A
DEFINE_*_BENCHMARKmacro that wrapsBENCHMARK_Pso each instantiation is a one-liner. - One registration line per
(device, dtype/layout/param)combination you want to measure.
A minimal example, modeled on benchmarks/src/roccv/bench_flip.cpp:
#include <core/hip_assert.h>
#include <core/image_format.hpp>
#include <core/tensor.hpp>
#include <op_flip.hpp>
#include <roccvbench/registry.hpp>
#include <roccvbench/utils.hpp>
#include "roccv_bench_helpers.hpp"
using namespace roccv;
template <eDeviceType DeviceType>
static roccvbench::BenchmarkResults RunFlipBenchmark(roccvbench::BenchmarkParamsList params) {
roccvbench::BenchmarkResults results;
// 1. Pull config + operator-specific params out of the params list.
ImageFormat inFormat = roccvbench::GetParamValue<ImageFormat>(params, "inFormat");
ImageFormat outFormat = roccvbench::GetParamValue<ImageFormat>(params, "outFormat");
int32_t flipCode = roccvbench::GetParamValue<int32_t>(params, "flipType");
int samples = roccvbench::GetParamValue<int>(params, "samples");
int width = roccvbench::GetParamValue<int>(params, "width");
int height = roccvbench::GetParamValue<int>(params, "height");
int runs = roccvbench::GetParamValue<int>(params, "runs");
int warmupRuns = roccvbench::GetParamValue<int>(params, "warmupRuns");
// 2. Allocate inputs/outputs on the target device.
Tensor input (Tensor::CalcRequirements(samples, {width, height}, inFormat, DeviceType));
Tensor output(Tensor::CalcRequirements(samples, {width, height}, outFormat, DeviceType));
// 3. Report memory footprint and fill the input with deterministic random data.
RegisterMemoryUsage(input, results.readMemoryBytes);
RegisterMemoryUsage(output, results.writtenMemoryBytes);
FillTensor(input);
// 4. Time the operator across runs + warmup.
Flip op;
hipStream_t stream;
HIP_VALIDATE_NO_ERRORS(hipStreamCreate(&stream));
roccvbench::RecordRuns<DeviceType>(stream, runs, warmupRuns, results.executionTimes,
[&]() { op(stream, input, output, flipCode, DeviceType); });
HIP_VALIDATE_NO_ERRORS(hipStreamDestroy(stream));
return results;
}
#define DEFINE_FLIP_BENCHMARK(name, device, inFormat, outFormat, flipType) \
BENCHMARK_P(Flip, name, \
BENCH_PARAMS(BENCH_PARAM("inFormat", inFormat), BENCH_PARAM("outFormat", outFormat), \
BENCH_PARAM("flipType", flipType))) { \
return RunFlipBenchmark<device>(params); \
}
// GPU
DEFINE_FLIP_BENCHMARK(GPU, eDeviceType::GPU, FMT_RGB8, FMT_RGB8, -1);
DEFINE_FLIP_BENCHMARK(GPU, eDeviceType::GPU, FMT_RGBA8, FMT_RGBA8, -1);
DEFINE_FLIP_BENCHMARK(GPU, eDeviceType::GPU, FMT_U8, FMT_U8, -1);
DEFINE_FLIP_BENCHMARK(GPU, eDeviceType::GPU, FMT_F32, FMT_F32, -1);
// CPU
DEFINE_FLIP_BENCHMARK(CPU, eDeviceType::CPU, FMT_RGB8, FMT_RGB8, -1);The registration macros live in benchmarks/roccvbench/include/roccvbench/registry.hpp:
| Macro | Purpose |
|---|---|
BENCHMARK(category, name) |
Define and register a parameter-less benchmark. The body must return BenchmarkResults. |
BENCHMARK_P(category, name, params) |
Same, but with a BenchmarkParamsList baked into the registration. |
BENCH_PARAMS(...) |
Wraps a comma-separated list of BENCH_PARAM(...) entries into a BenchmarkParamsList. |
BENCH_PARAM(key, value) |
One typed parameter; the stringified value is used as the display label. |
BENCH_PARAM_STR(key, value, strValue) |
Same, but lets you override the display string when the macro stringification is unreadable. |
Category groups related benchmarks (typically the operator's CamelCase name, e.g. Flip, Resize, CvtColor). The --select / --exclude CLI flags operate on categories. Name is the per-instantiation label (GPU, CPU, GPU_U8, etc.) — use it to distinguish variants when one operator has many instantiations.
Two distinct sources feed params:
-
Suite-wide config, injected automatically by the harness from
benchmarks/config.json:samples,width,height,runs,warmupRuns. Read these in every benchmark — never hardcode image dimensions or run counts. -
Per-benchmark parameters, set when you register the unit via
BENCH_PARAMS(...). These are how you sweep the operator's own knobs (interpolation type, color conversion code, format pair, flip axis, …).
Inside the runner, read each parameter with roccvbench::GetParamValue<T>(params, "key"). The type must match what you stored — int, int32_t, ImageFormat, eInterpolationType, eColorConversionCode, etc. A mismatched type throws via std::any_cast.
From benchmarks/src/roccv/roccv_bench_helpers.hpp:
-
FillTensor(const Tensor&)— fills with deterministic random data using the suite-wide seed (kBenchSeedinbenchmarks/roccvbench/include/roccvbench/utils.hpp). Always use this rather than rolling your own; data-dependent operators (e.g. threshold) need byte-identical inputs across sessions to be comparable. -
RegisterMemoryUsage(const Tensor&, size_t&)— accumulates the tensor's byte size into eitherresults.readMemoryBytes(inputs) orresults.writtenMemoryBytes(outputs). The analysis pipeline uses these to compute effective bandwidth.
From benchmarks/roccvbench/include/roccvbench/utils.hpp:
-
RecordRuns<DeviceType>(stream, runs, warmupRuns, executionTimes, fn)— runsfnruns + warmupRunstimes, discards the warmup runs, appends each timed run's duration (in seconds) toexecutionTimes. Uses HIP events for GPU andstd::chrono::steady_clockfor CPU, so the GPU timing is on-device kernel time only. -
RecordRunsCpu(runs, warmupRuns, executionTimes, fn)— CPU-only convenience overload for paths that don't have a HIP stream (e.g. OpenCV reference benches).
The instantiation list is where you decide what gets measured. A few rules of thumb:
-
Cover each
(dtype, channels)dispatcher branch the operator actually supports, at least on GPU — that's what catches a kernel regression in one type but not another. - Sweep the operator's own knobs when they meaningfully change the kernel path (interpolation type, border type, conversion code). Don't sweep every legal combination; pick the representative ones.
-
Don't sweep
samples/width/heightper-bench — those are driven byconfig.json, which the harness iterates over for every registered unit.
benchmarks/src/roccv/bench_resize.cpp and benchmarks/src/roccv/bench_cvt_color.cpp are good references for operators with more than one extra parameter.
After dropping the file in and reconfiguring CMake:
cmake --build . --parallel
./bin/roccv_bench --list # confirm the category appears
./bin/roccv_bench --select <Category> # run just yoursA successful list output and a non-empty result row per instantiation is enough to confirm the unit is wired up. For interpreting and validating the actual numbers, hand off to benchmarks/README.md and benchmarks/analyze_results.py.
- Home
- Contributing
-
Architecture
- Overview
- Containers
- Kernel Wrappers
- Operator Flow