mirror of
https://github.com/NVIDIA/nvbench.git
synced 2026-06-29 18:57:44 +00:00
d63a2761eb
* Add type annotations for future functionality
```python
class Timer:
def start(self) -> None: ...
def stop(self) -> None: ...
```
and overloaded `State.exec` so:
- normal mode accepts `Callable[[Launch], None]`
- `timer=True` accepts `Callable[[Launch, Timer], None]`
No implementation yet. Type annotation checked with
```
(py313) :~/repos/nvbench/python$ python -m mypy --ignore-missing-imports /tmp/check_timer.py
/tmp/check_timer.py:24: error: No overload variant of "exec" of "State" matches argument types "Callable[[Launch], None]", "bool" [call-overload]
/tmp/check_timer.py:24: note: Possible overload variants:
/tmp/check_timer.py:24: note: def exec(self, Callable[[Launch], None], /, *, batched: bool | None = ..., sync: bool | None = ..., timer: Literal[False] = ...) -> None
/tmp/check_timer.py:24: note: def exec(self, Callable[[Launch, Timer], None], /, *, timer: Literal[True], sync: bool | None = ...) -> None
/tmp/check_timer.py:25: error: Argument 1 to "exec" of "State" has incompatible type "Callable[[Launch, Timer], None]"; expected "Callable[[Launch], None]" [arg-type]
/tmp/check_timer.py:26: error: No overload variant of "exec" of "State" matches argument types "Callable[[Launch, int], None]", "bool" [call-overload]
/tmp/check_timer.py:26: note: Possible overload variants:
/tmp/check_timer.py:26: note: def exec(self, Callable[[Launch], None], /, *, batched: bool | None = ..., sync: bool | None = ..., timer: Literal[False] = ...) -> None
/tmp/check_timer.py:26: note: def exec(self, Callable[[Launch, Timer], None], /, *, timer: Literal[True], sync: bool | None = ...) -> None
Found 3 errors in 1 file (checked 1 source file)
(py313) :~/repos/nvbench/python$ nl -ba /tmp/check_timer.py
1 # /tmp/check_nvbench_timer.py
2 import cuda.bench as bench
3
4 def normal_ok(launch: bench.Launch) -> None:
5 pass
6
7 def timer_ok(launch: bench.Launch, timer: bench.Timer) -> None:
8 timer.start()
9 timer.stop()
10
11 def missing_timer(launch: bench.Launch) -> None:
12 pass
13
14 def extra_timer(launch: bench.Launch, timer: bench.Timer) -> None:
15 pass
16
17 def wrong_timer_type(launch: bench.Launch, timer: int) -> None:
18 pass
19
20 def state_bench(state: bench.State) -> None:
21 state.exec(normal_ok)
22 state.exec(normal_ok, timer=False)
23 state.exec(timer_ok, timer=True)
24 state.exec(missing_timer, timer=True) # should fail
25 state.exec(extra_timer) # should fail
26 state.exec(wrong_timer_type, timer=True) # should fail
```
* Implement cuda.bench.Timer object
The Timer class is not user-constructible. It exposes two nullary
methods timer.start() and timer.stop().
The instance of Timer class would be provided to launchable object
passed to State.exec with timer=True.
* Implement support for State.exec( launch_fn, timer=True)
* Change type annotation for batch to default to None
None is interpreted as `not timer`, i.e., it effectively
defaults to True (as before) for usage without timer set,
but starts defaulting to `False` is `timer=True` is set.
The batched keyword type is `bool | None`.
* Implement default batched=None behavior
API allows one to specify all 3 keywords, sync, batched,
and timer. batched is None by default, run-time interpreted
as `(not timer)`.
* Update tests for new behavior of batched/time combination
* Add python/examples/exec_tag_timer.py
* Expand Timer class and methods docstrings
* Reworked python/example/exec_tag_timer.py to align with C++ example.
* Replace ::cuda::std::name with cuda::std::name
* Resolve review feedback
89 lines
2.7 KiB
Python
89 lines
2.7 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 cuda.bench as bench
|
|
import cuda.cccl.headers as headers
|
|
import cuda.core as core
|
|
|
|
|
|
def as_core_Stream(cs: bench.CudaStream) -> core.Stream:
|
|
"Create view into native stream provided by NVBench"
|
|
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.float64("Duration", [1e-4 + k * 0.25e-3 for k in range(5)])
|
|
@bench.axis.string("Kramble", ["Foo", "Bar", "Baz"])
|
|
def runtime_skip(state: bench.State):
|
|
duration = state.get_float64("Duration")
|
|
kramble = state.get_string("Kramble")
|
|
|
|
# Skip Baz benchmarks with 0.8 ms duration
|
|
if kramble == "Baz" and duration < 0.8e-3:
|
|
state.skip("Short 'Baz' benchmarks are skipped")
|
|
return
|
|
|
|
# Skip Foo benchmark with > 0.3 ms duration
|
|
if kramble == "Foo" and duration > 0.3e-3:
|
|
state.skip("Long 'Foo' benchmarks are skipped")
|
|
return
|
|
|
|
krn = make_sleep_kernel()
|
|
launch_cfg = core.LaunchConfig(grid=1, block=1, shmem_size=0)
|
|
|
|
def launcher(launch: bench.Launch):
|
|
s = as_core_Stream(launch.get_stream())
|
|
core.launch(s, launch_cfg, krn, duration)
|
|
|
|
state.exec(launcher)
|
|
|
|
|
|
if __name__ == "__main__":
|
|
bench.run_all_benchmarks(sys.argv)
|