nvbench: CUDA kernel benchmarking with parameter axes, clock locking and throughput math

Built for one kernel at a time, and says so up front

The README states the scope plainly, and it is worth reading before anything else. 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 across multiple applications, the README points to the NVIDIA Nsight tools as the appropriate choice.

That boundary is the honest core of the project. A lot of GPU benchmarking tools blur the line between measuring a kernel and measuring an application, and then produce numbers that are hard to compare. nvbench refuses that blur. One benchmark, one region, one number.

The other half of the scope statement covers CPU work. The library provides CPU-only benchmarking facilities for non-trivial CPU workloads, but the README says it is not optimized for CPU microbenchmarks and suggests Google Benchmark for high resolution CPU benchmarking. It also notes that the CPU-only mode is not suitable for microbenchmarking.

The README opens with a warning that this project is a work in progress and that everything is subject to change. For a library at version 0.x with a separate Python package at 0.3.0, that framing is accurate rather than hedging, and it should shape how you pin versions.

The audience is CUDA kernel developers, mostly. The topics listed on the repository are benchmark, cuda, cuda-kernels, gpu, kernel-benchmark, nvidia and performance, which is as narrow and as useful a topic list as you will find.

Parameter axes are the core idea, not a convenience

Most of the design effort in nvbench goes into the axis system. An axis is a parameter of the benchmark, and the library explores the configuration space defined by those axes. The README describes the system as powerful and flexible, and notes that parameters may be dynamic numbers or strings, or static types, which matters for axes that are compile-time constants such as enum values.

The distinction between dynamic and static axes is the part with real consequences. A dynamic axis can be swept at runtime, meaning one compiled binary tests many parameter values. A static axis changes what the compiler sees, so the sweep happens across builds, and the examples directory includes a dedicated file for enums and compile-time-constant integral axes.

The examples directory shows the range of what the system covers: runtime and compile-time parameter sweeps in `axes.cu`, enums and compile-time constants in `enums.cu`, CPU-only benchmarking in `cpu_only.cu`, throughput reporting in `throughput.cu`, skipping configurations in `skip.cu`, stream selection in `stream.cu`, adding or hiding markdown summary columns in `summaries.cu`, and the two execution tags in `exec_tag_sync.cu` and `exec_tag_timer.cu`. There is also `auto_throughput.cu` and `custom_criterion.cu`, which go beyond what the README lists.

The payoff of a sweep-based system is regression detection rather than a single best number. When every configuration is measured and printed in a table, a change that makes one configuration slower is visible in the diff of the output, and the skip mechanism exists so a configuration known to be invalid can be excluded rather than producing a spurious failure.

Three measurement modes, and what cold actually means

The measurement types are where a benchmark earns or loses trust, and nvbench separates three of them.

Cold measurement runs the benchmark once per sample with a clean device L2 cache, and reports both GPU and CPU times. The cold distinction is the important one: a kernel that reads a large working set will look faster on a warm cache than on a cold one, and a benchmark that does not control for it produces numbers that depend on what ran before it.

Batch measurement executes the benchmark multiple times back to back and records total time, reporting the average execution time as total time divided by the number of executions. This is the mode you want when the kernel is short enough that launch overhead dominates the single-shot measurement.

CPU-only measurement records host-side execution time for a non-GPU benchmark, with the README noting it is not suitable for microbenchmarking.

There is also a manual timer mode, described as explicitly starting and stopping timing inside a benchmark implementation, which corresponds to the `nvbench::exec_tag::timer` tag in the examples. The reason to reach for it is that automatic timing includes overhead the author does not want in the number, and manual timing lets the author exclude it. The companion `nvbench::exec_tag::sync` tag exists for benchmarks that synchronise CUDA devices, which matters because an unsynchronised benchmark measures launch rather than completion.

Both tags are worth understanding before you trust a result. A benchmark that measures without synchronising can report the time to enqueue work rather than the time to run it.

The command line can redefine the sweep and lock the clocks

Each benchmark executable NVBench produces carries its own command line options for configuring execution at runtime, and the README highlights two of them as distinctive.

The first is redefinition of parameter axes. The axes compiled into the binary can be replaced or narrowed at invocation time, so one binary built with a wide sweep can be run with a narrow subset without recompiling. The second is CUDA device selection, which makes it practical to run the same binary across several GPUs in a machine.

The third is locking GPU clocks, supported on Volta and later. This is the option that separates a benchmark harness from a homegrown timer. Unlocked clocks drift with temperature and power state, so the same kernel can measure differently on two runs for reasons that have nothing to do with the code. The sample output in the README shows the noise tracking that results from this: the run reports a measurement timing out at 15.00s while over a noise threshold of 1.26% against 0.50%.

The other options cover changing output formats and other execution settings. The supported formats are markdown, which is the default, and CSV. Markdown being the default is a deliberate choice for a tool aimed at regression testing, because a markdown table is what a reviewer reads in a pull request, and the summaries example shows how columns can be added or hidden to shape that table.

The linked talk on CUDA kernel benchmarking covers the problems the library is designed to solve, which is a reasonable place to go if the docs feel thin on any of this.

A minimal benchmark is four lines of CUDA

The Getting Started section shows that the entry cost is small. A benchmark is a function taking an `nvbench::state`, which executes the work through a launch object, plus a macro that registers it:

void my_benchmark(nvbench::state& state) {
  state.exec([](nvbench::launch& launch) {
    my_kernel<<<num_blocks, 256, 0, launch.get_stream()>>>();
  });
}
NVBENCH_BENCH(my_benchmark);

The structure encodes the measurement discipline. The kernel launch is passed the stream from the launch object rather than the default stream, so the benchmark participates in the harness stream setup instead of running on stream zero behind the harness's back. The kernel is launched inside `state.exec`, which is where the timing, the repetition count and the cache handling live, so a kernel launched outside that call is not measured at all.

Building the examples needs the examples flag turned on and the CUDA architecture set for the GPU you will run on:

mkdir -p build
cd build
cmake -DNVBench_ENABLE_EXAMPLES=ON -DCMAKE_CUDA_ARCHITECTURES=70 .. && make

The README warns to set `CMAKE_CUDA_ARCHITECTURE` based on the GPU, and the example value is 70. The binaries land in `build/bin` with the prefix `nvbench.example`.

For the library itself, the supported toolchain is pinned in the README: CMake above 3.30.4, CUDA Toolkit with nvcc 12.0 and above, g++ from 7 through 14, clang++ from 14 through 19, and headers tested from C++17 through C++20. The library is C++17. That is a narrow enough matrix to be worth checking before you start, since a compiler outside those ranges will not have been tested.

Throughput turns a time into a rate, and the cache and cold start explain the rest

A time in microseconds is not very meaningful on its own, and nvbench computes throughput measures alongside it. The README lists item throughput in elements per second, and global memory bandwidth usage in bytes per second and as a percentage of the device peak bandwidth.

Expressing bandwidth as a percentage of peak is the more useful of the two, because it tells you how close a kernel is to the memory system limit. A kernel at a small fraction of peak bandwidth is usually limited by something other than memory, and the number narrows down where to look. The example output shows this in the `GlobalMem BW` and `BWPeak` columns, with the peak value as a percentage.

The measurement itself comes with the noise analysis, which is what makes the numbers defensible. The sample run reports a cold GPU time of 262.392 us at 1.26% noise, a batch GPU time of 261.963 us, and a batch count of 27394x. The repeated execution count is reported alongside, so you can see how many samples backed the figure. The log line for the run also shows the timeout behaviour, where a measurement exceeded its 15 second budget while still being over the noise threshold.

Two features of the sample output deserve a second look. The device section prints the hardware facts a benchmark needs for interpretation: SM count, SM clock, total and free global memory, peak memory bus bandwidth, L2 cache size, shared memory per SM and per block, maximum active blocks and threads per SM, register availability and ECC status. And the table includes samples, CPU time, GPU time, noise for each, element throughput, bandwidth, bandwidth peak percentage and the batch columns.

For a kernel whose performance depends on occupancy or cache behaviour, having those device limits printed next to the measurement means you can tell whether a result is explained by the hardware or by the code.

A separate Python package, and an Apache 2.0 licence with LLVM exceptions

The repository is not only C++. There is a `python/` directory and a `pyproject.toml`, and the release history shows a Python library with its own tags: python-0.3.0 on 2026-08-05, python-0.2.0 on 2026-02-04 and python-0.1.0 on 2026-01-30. The repository root also carries `requirements.txt` listing colorama, matplotlib, numpy, pandas, seaborn and tabulate, which is a plotting and tabular output stack rather than a benchmark harness stack, consistent with a tool for turning benchmark results into charts.

The root `pyproject.toml` is configured for ruff with a py310 target, and it configures codespell for pre-commit. It sits alongside a `CMakeLists.txt`, `CMakePresets.json` and a `cmake/` directory, so the repository holds a C++ library, a Python layer and a documentation toolchain at once. There is also an `exec/` directory, a `testing/` directory and a `ci/` directory for continuous integration, plus a `docs/` directory holding the markdown pages the README links to, including the benchmarks, CLI help and axis specification pages.

The licence is Apache 2.0 with LLVM exceptions, as the README's licence section states, and the repository licence field records Apache-2.0. The LLVM exceptions matter in practice: they let the code be included in projects that are distributed under the same terms as LLVM, which is what makes it usable inside compiler and GPU toolchain codebases without a relicensing negotiation. Anyone redistributing it should read the LICENSE file rather than relying on the badge.

The last push was on 2026-10-14, so the code is moving. The project also has a CODE_OF_CONDUCT.md, an AGENTS.md file, a .pre-commit-config.yaml and a .devcontainer directory, which is the usual sign of a project with a real contribution process rather than a single maintainer's script.

Editorial conclusion

nvbench suits a CUDA developer who needs to compare kernel configurations and track them in CI, since the axis system and the markdown output are built for exactly that, and it is the wrong tool for end to end application profiling, where the README points to the NVIDIA Nsight tools, or for CPU microbenchmarks, where it recommends Google Benchmark instead. Verify three things before adopting it: that your toolchain is inside the supported matrix, CMake above 3.30.4 with CUDA Toolkit 12.0 or newer and g++ 7 through 14 or clang++ 14 through 19, that you set `CMAKE_CUDA_ARCHITECTURES` to match the GPU you run on, and that you read the project status line, because the README describes NVBench as a work in progress where everything is subject to change.

Frequently asked questions

Official sources

1. Issues
2. License: Apache-2.0
3. NVIDIA/nvbench on GitHub
4. README
5. Releases