mirror of
https://github.com/NVIDIA/nvbench.git
synced 2026-07-13 02:27:22 +00:00
d483a5bc73
* Replace forwarding with semantically more accurate std::move Also add comment within percentile_rank to document precondition on input values checked with assert statement. Also, sharpened the comment around percentile_rank function * Duplicate-heavy boundary test is added Prepare duplicate heavy input and check sort-based quartile computation result with selection-based one. std::nth_element only guarantees that the nth element is the value that would appear there in sorted order; it does not fully sort equal partitions. Bugs in the selection implementation, especially when selecting Q1 from the left half and Q3 from the right half after selecting the median, are more likely to show up when many samples equal the quartile values. * Generate cold summaries only if some accepted samples have been accumulated Cold measurement can discard throttled trials before incrementing the accepted sample count, then stop on timeout with zero recorded samples. In that case, only emit the sample-size summary and skip derived timing, bandwidth, clock, and bulk summaries that require accepted samples. This avoids divide-by-zero mean calculations and quartile/IQR computation over empty sample vectors. Keep timeout diagnostics reachable for zero-sample runs and add an explicit warning when no accepted cold samples were recorded. Factor timeout warning emission into a private helper so the zero-sample and normal paths share the same diagnostic logic. Suppress low-sample relative stdev noise Add a statistics helper that returns no relative standard-deviation noise until there are enough samples for a meaningful estimate. Use it for cold CPU/GPU and CPU-only summaries so the low-sample +inf stdev sentinel is not published as real relative noise or used for max-noise timeout warnings. Add statistics coverage for suppressing the low-sample sentinel and computing relative stdev noise once the sample threshold is reached. compute_standard_deviation_noise return nullopt if standard deviation is not finite Test verify that noise is nullopt when not enough samples are accumulated Added statistics::has_enough_samples_for_noise_estimate(...) Used it in standard_deviation, compute_standard_deviation_noise, compute_robust_noise. Added timeout diagnostics in cold and CPU-only paths. if max-noise is configured and the run timed out before enough samples exist to estimate noise, the log now says that explicitly, otherwise the existing “over noise threshold” warning remains unchanged. Added a statistics test assertion for the new sample-count predicate. * Test quartile values across selection threshold Add fixed expected-value assertions for quartile tests around the sort/selection switch point, including duplicate-heavy inputs. This keeps the tests from only proving that both implementations agree with each other. * Add NaN guards to percentiles and quartiles computation routines * Add tests for handling of NaNs in quartile routine inputs * Add comment re magic sort/select threshold value * Refactor logic of emitting warnings between cold and cpu-only measures Introduce new header file with inline implementation. Use it from measure_cold.cuh and measure_cpu_only.cxx * Add static assertion that ValueType is a floating-point type * Check consistency of sort- vs. select-based quartiles using threshold constant Expose quartile threshold value, use it in testing to test around that value. * Collapsed two branches with identical bodies * test_compute_standard_deviation_noise exercises other invalid inputs * Preserve stdev noise summaries for low sample counts Keep legacy stdev/relative summary tags present even when too few samples are available to compute a meaningful standard-deviation noise estimate. Use the standard-deviation unavailable sentinel for those values so existing summary consumers continue to see the expected tags. Factor the sentinel into the statistics helpers and use it from both standard_deviation() and stdev_noise_or_sentinel(), keeping the schema compatibility behavior explicit and tested. * timeout_warnings now treats engaged NaN and negative stdev noise as unavailable Add a focused test target, nvbench.test.measure_timeout_warnings, covering: - NaN stdev noise -> “unable to estimate noise” - negative stdev noise -> “unable to estimate noise” - +inf stdev noise -> “over noise threshold” * Test nullopt explicitly in warning check test check_noise_warning() now takes std::optional<nvbench::float64_t>, matching the production helper, and the test now covers std::nullopt explicitly in addition to NaN, negative, and +inf. * Tighten statistics and timeout warning tests Document that percentile helpers return quiet NaNs for NaN-containing inputs. Make quartile expected-value tests compute ranks from the documented round(p / 100 * (n - 1)) rule instead of reusing statistics::percentile_rank(), so rank regressions are caught independently. Extend timeout-warning coverage to exercise the too-few-samples max-noise path in addition to unavailable, invalid, and infinite stdev-noise inputs. * Test for timeout warnings for min-samples and min-time
Overview
This project is a work-in-progress. Everything is subject to change.
NVBench is a C++17 library designed to simplify CUDA kernel benchmarking. It features:
- Parameter sweeps: a powerful and flexible "axis" system explores a kernel's configuration space. Parameters may be dynamic numbers/strings or static types.
- Runtime customization: A rich command-line interface allows redefinition of parameter axes, CUDA device selection, locking GPU clocks (Volta+), changing output formats, and more.
- Throughput calculations: Compute
and report:
- Item throughput (elements/second)
- Global memory bandwidth usage (bytes/second and per-device %-of-peak-bw)
- Multiple output formats: Currently supports markdown (default) and CSV output.
- Manual timer mode: (optional) Explicitly start/stop timing in a benchmark implementation.
- Multiple measurement types:
- Cold Measurements:
- Each sample runs the benchmark once with a clean device L2 cache.
- GPU and CPU times are reported.
- Batch Measurements:
- Executes the benchmark multiple times back-to-back and records total time.
- Reports the average execution time (total time / number of executions).
- CPU-only Measurements
- Measures the host-side execution time of a non-GPU benchmark.
- Not suitable for microbenchmarking.
- Cold Measurements:
Check out this talk for an overview of the challenges inherent to CUDA kernel benchmarking and how NVBench solves them for you!
Supported Compilers and Tools
- CMake > 3.30.4
- CUDA Toolkit + nvcc: 12.0 and above
- g++: 7 -> 14
- clang++: 14 -> 19
- Headers are tested with C++17 -> C++20.
Getting Started
Minimal Benchmark
A basic kernel benchmark can be created with just a few lines of CUDA C++:
void my_benchmark(nvbench::state& state) {
state.exec([](nvbench::launch& launch) {
my_kernel<<<num_blocks, 256, 0, launch.get_stream()>>>();
});
}
NVBENCH_BENCH(my_benchmark);
See Benchmarks for information on customizing benchmarks and implementing parameter sweeps.
Command Line Interface
Each benchmark executable produced by NVBench provides a rich set of command-line options for configuring benchmark execution at runtime. See the CLI overview and CLI axis specification for more information.
Examples
This repository provides a number of examples that demonstrate various NVBench features and usecases:
- Runtime and compile-time parameter sweeps
- CPU-only benchmarking
- Enums and compile-time-constant-integral parameter axes
- Reporting item/sec and byte/sec throughput statistics
- Skipping benchmark configurations
- Benchmarking on a specific stream
- Adding / hiding columns (summaries) in markdown output
- Benchmarks that sync CUDA devices:
nvbench::exec_tag::sync - Manual timing:
nvbench::exec_tag::timer
Building Examples
To build the examples:
mkdir -p build
cd build
cmake -DNVBench_ENABLE_EXAMPLES=ON -DCMAKE_CUDA_ARCHITECTURES=70 .. && make
Be sure to set CMAKE_CUDA_ARCHITECTURE based on the GPU you are running on.
Examples are built by default into build/bin and are prefixed with nvbench.example.
Example output from `nvbench.example.throughput`
# Devices
## [0] `Quadro GV100`
* SM Version: 700 (PTX Version: 700)
* Number of SMs: 80
* SM Default Clock Rate: 1627 MHz
* Global Memory: 32163 MiB Free / 32508 MiB Total
* Global Memory Bus Peak: 870 GiB/sec (4096-bit DDR @850MHz)
* Max Shared Memory: 96 KiB/SM, 48 KiB/Block
* L2 Cache Size: 6144 KiB
* Maximum Active Blocks: 32/SM
* Maximum Active Threads: 2048/SM, 1024/Block
* Available Registers: 65536/SM, 65536/Block
* ECC Enabled: No
# Log
Run: throughput_bench [Device=0]
Warn: Current measurement timed out (15.00s) while over noise threshold (1.26% > 0.50%)
Pass: Cold: 0.262392ms GPU, 0.267860ms CPU, 7.19s total GPU, 27393x
Pass: Batch: 0.261963ms GPU, 7.18s total GPU, 27394x
# Benchmark Results
## throughput_bench
### [0] Quadro GV100
| NumElements | DataSize | Samples | CPU Time | Noise | GPU Time | Noise | Elem/s | GlobalMem BW | BWPeak | Batch GPU | Batch |
|-------------|------------|---------|------------|-------|------------|-------|---------|---------------|--------|------------|--------|
| 16777216 | 64.000 MiB | 27393x | 267.860 us | 1.25% | 262.392 us | 1.26% | 63.940G | 476.387 GiB/s | 58.77% | 261.963 us | 27394x |
Demo Project
To get started using NVBench with your own kernels, consider trying out the NVBench Demo Project.
nvbench_demo provides a simple CMake project that uses NVBench to build an
example benchmark. It's a great way to experiment with the library without a lot
of investment.
Contributing
Contributions are welcome!
For current issues, see the issue board. Issues labeled with 
Tests
To build nvbench tests:
mkdir -p build
cd build
cmake -DNVBench_ENABLE_TESTING=ON .. && make
Tests are built by default into build/bin and prefixed with nvbench.test.
To run all tests:
make test
or
ctest
License
NVBench is released under the Apache 2.0 License with LLVM exceptions. See LICENSE.
Scope and Related Projects
NVBench will measure the CPU and CUDA GPU execution time of a single host-side critical region per benchmark. It is intended for regression testing and parameter tuning of individual kernels. For in-depth analysis of end-to-end performance of multiple applications, the NVIDIA Nsight tools are more appropriate.
NVBench is focused on evaluating the performance of CUDA kernels. It also provides CPU-only benchmarking facilities intended for non-trivial CPU workloads, but is not optimized for CPU microbenchmarks. This may change in the future, but for now, consider using Google Benchmark for high resolution CPU benchmarks.