mirror of
https://github.com/NVIDIA/nvbench.git
synced 2026-06-29 02:37:36 +00:00
44ec7de6bd
* Add decorators for registering benchmarks and adding axis
cuda.bench.register(fn) continues returning Benchmark, and supports
legacy use.
New signature added:
cuda.bench.register():
Returns a decorator
```
@bench.register()
@bench.axis.float64("Duration (s)", [7e-5, 1e-4, 5e-4])
@bench.option.min_samples(120)
def single_float64_axis(state: bench.State):
...
```
* Remove example/auto_throughput.py
The C++ counterpart's purpose is to demonstrate use of CUPTI
metrics, but these are not supported in Python bindings, so
this example is a duplicate of example/throughput.py
* Add wrong decorator order test for bench.axis.*
* Strengthen type annotation for register function
Acting on code rabbit nit-pick require that function being
registered take cuda.bench.State object as an argument.
Verified the fix as
```
(py313) :~/repos/nvbench/python$ python -m mypy --ignore-missing-import /tmp/t.py
/tmp/t.py:8: error: Argument 1 has incompatible type "Callable[[], None]"; expected "Callable[[State], None]" [arg-type]
Found 1 error in 1 file (checked 1 source file)
(py313) :~/repos/nvbench/python$ nl -ba /tmp/t.py
1 # /tmp/check_nvbench_register.py
2 import cuda.bench as bench
3
4 @bench.register()
5 def good(state: bench.State) -> None:
6 pass
7
8 @bench.register()
9 def bad() -> None:
10 pass
```
* Replace use of global variable with thread-safe lru_cache
This improves thread-safety of module initialization.
* Abide by RUF005 linting rule
* Expand docstrings regarding cuda.bench.register() decorator
It explains to the user what the decorator does and provides
a concise usage example.
* Sharpen wording on exception maybe-thrown by decorator
93 lines
2.6 KiB
Python
93 lines
2.6 KiB
Python
# Copyright 2025-2026 NVIDIA Corporation
|
|
#
|
|
# Licensed under the Apache License, Version 2.0 with the LLVM exception
|
|
# (the "License"); you may not use this file except in compliance with
|
|
# the License.
|
|
#
|
|
# You may obtain a copy of the License at
|
|
#
|
|
# http://llvm.org/foundation/relicensing/LICENSE.txt
|
|
#
|
|
# Unless required by applicable law or agreed to in writing, software
|
|
# distributed under the License is distributed on an "AS IS" BASIS,
|
|
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
|
# See the License for the specific language governing permissions and
|
|
# limitations under the License.
|
|
|
|
import sys
|
|
import time
|
|
|
|
import cuda.bench as bench
|
|
import cuda.cccl.headers as headers
|
|
import cuda.core as core
|
|
|
|
host_sleep_duration = 0.1
|
|
|
|
|
|
@bench.register()
|
|
@bench.option.is_cpu_only()
|
|
def cpu_only_sleep_bench(state: bench.State) -> None:
|
|
def launcher(launch: bench.Launch):
|
|
time.sleep(host_sleep_duration)
|
|
|
|
state.exec(launcher)
|
|
|
|
|
|
def as_core_Stream(cs: bench.CudaStream) -> core.Stream:
|
|
return core.Stream.from_handle(cs.addressof())
|
|
|
|
|
|
def make_sleep_kernel():
|
|
"""JITs sleep_kernel(seconds)"""
|
|
src = r"""
|
|
#include <cuda/std/cstdint>
|
|
#include <cuda/std/chrono>
|
|
|
|
// Each launched thread just sleeps for `seconds`.
|
|
__global__ void sleep_kernel(double seconds) {
|
|
namespace chrono = ::cuda::std::chrono;
|
|
using hr_clock = chrono::high_resolution_clock;
|
|
|
|
auto duration = static_cast<cuda::std::int64_t>(seconds * 1e9);
|
|
const auto ns = chrono::nanoseconds(duration);
|
|
|
|
const auto start = hr_clock::now();
|
|
const auto finish = start + ns;
|
|
|
|
auto now = hr_clock::now();
|
|
while (now < finish)
|
|
{
|
|
now = hr_clock::now();
|
|
}
|
|
}
|
|
"""
|
|
incl = headers.get_include_paths()
|
|
opts = core.ProgramOptions(include_path=str(incl.libcudacxx))
|
|
prog = core.Program(src, code_type="c++", options=opts)
|
|
mod = prog.compile("cubin", name_expressions=("sleep_kernel",))
|
|
return mod.get_kernel("sleep_kernel")
|
|
|
|
|
|
@bench.register()
|
|
@bench.axis.string("Sync", ["Do not sync", "Do sync"])
|
|
def mixed_sleep_bench(state: bench.State) -> None:
|
|
sync = state.get_string("Sync")
|
|
sync_flag = sync == "Do sync"
|
|
|
|
gpu_sleep_dur = 225e-3
|
|
krn = make_sleep_kernel()
|
|
launch_config = core.LaunchConfig(grid=1, block=1, shmem_size=0)
|
|
|
|
def launcher(launch: bench.Launch):
|
|
# host overhead
|
|
time.sleep(host_sleep_duration)
|
|
# GPU computation
|
|
s = as_core_Stream(launch.get_stream())
|
|
core.launch(s, launch_config, krn, gpu_sleep_dur)
|
|
|
|
state.exec(launcher, sync=sync_flag)
|
|
|
|
|
|
if __name__ == "__main__":
|
|
bench.run_all_benchmarks(sys.argv)
|