diff --git a/python/cuda/bench/__init__.py b/python/cuda/bench/__init__.py index 119f994..6628e32 100644 --- a/python/cuda/bench/__init__.py +++ b/python/cuda/bench/__init__.py @@ -42,6 +42,7 @@ _NVBENCH_EXPORTS = ( "Launch", "NVBenchRuntimeError", "State", + "Timer", "run_all_benchmarks", ) diff --git a/python/cuda/bench/__init__.pyi b/python/cuda/bench/__init__.pyi index dd69e29..f9d4964 100644 --- a/python/cuda/bench/__init__.pyi +++ b/python/cuda/bench/__init__.pyi @@ -28,6 +28,7 @@ from collections.abc import Callable, Sequence from typing import ( Any, + Literal, Optional, Self, SupportsFloat, @@ -67,6 +68,10 @@ class Benchmark: class Launch: def get_stream(self) -> CudaStream: ... +class Timer: + def start(self) -> None: ... + def stop(self) -> None: ... + class State: def has_device(self) -> bool: ... def has_printers(self) -> bool: ... @@ -107,14 +112,26 @@ class State: def set_timeout(self, duration: SupportsFloat) -> None: ... def get_blocking_kernel_timeout(self) -> float: ... def set_blocking_kernel_timeout(self, duration: SupportsFloat) -> None: ... + @overload def exec( self, fn: Callable[[Launch], None], /, *, - batched: Optional[bool] = True, - sync: Optional[bool] = False, - ): ... + batched: Optional[bool] = None, + sync: bool = False, + timer: Literal[False] = False, + ) -> None: ... + @overload + def exec( + self, + fn: Callable[[Launch, Timer], None], + /, + *, + timer: Literal[True], + batched: Literal[False] | None = None, + sync: bool = False, + ) -> None: ... def get_short_description(self) -> str: ... def add_summary( self, column_name: str, value: Union[SupportsInt, SupportsFloat, str] diff --git a/python/examples/axes.py b/python/examples/axes.py index 6afcedb..17ff231 100644 --- a/python/examples/axes.py +++ b/python/examples/axes.py @@ -35,7 +35,7 @@ def make_sleep_kernel(): // Each launched thread just sleeps for `seconds`. __global__ void sleep_kernel(double seconds) { - namespace chrono = ::cuda::std::chrono; + namespace chrono = cuda::std::chrono; using hr_clock = chrono::high_resolution_clock; auto duration = static_cast(seconds * 1e9); @@ -101,7 +101,7 @@ def make_copy_kernel(in_type: Optional[str] = None, out_type: Optional[str] = No * Naive copy of `n` values from `in` -> `out`. */ template -__global__ void copy_kernel(const T *in, U *out, ::cuda::std::size_t n) +__global__ void copy_kernel(const T *in, U *out, cuda::std::size_t n) { const auto init = blockIdx.x * blockDim.x + threadIdx.x; const auto step = blockDim.x * gridDim.x; @@ -116,9 +116,9 @@ __global__ void copy_kernel(const T *in, U *out, ::cuda::std::size_t n) opts = core.ProgramOptions(include_path=str(incl.libcudacxx)) prog = core.Program(src, code_type="c++", options=opts) if in_type is None: - in_type = "::cuda::std::int32_t" + in_type = "cuda::std::int32_t" if out_type is None: - out_type = "::cuda::std::int32_t" + out_type = "cuda::std::int32_t" instance_name = f"copy_kernel<{in_type}, {out_type}>" mod = prog.compile("cubin", name_expressions=(instance_name,)) return mod.get_kernel(instance_name) diff --git a/python/examples/cpu_activity.py b/python/examples/cpu_activity.py index c724501..b7c159f 100644 --- a/python/examples/cpu_activity.py +++ b/python/examples/cpu_activity.py @@ -45,7 +45,7 @@ def make_sleep_kernel(): // Each launched thread just sleeps for `seconds`. __global__ void sleep_kernel(double seconds) { - namespace chrono = ::cuda::std::chrono; + namespace chrono = cuda::std::chrono; using hr_clock = chrono::high_resolution_clock; auto duration = static_cast(seconds * 1e9); diff --git a/python/examples/exec_tag_sync.py b/python/examples/exec_tag_sync.py index cf4ec29..c008ea7 100644 --- a/python/examples/exec_tag_sync.py +++ b/python/examples/exec_tag_sync.py @@ -36,7 +36,7 @@ def make_fill_kernel(data_type: Optional[str] = None): * Naive setting of values in buffer */ template -__global__ void fill_kernel(T *buf, T v, ::cuda::std::size_t n) +__global__ void fill_kernel(T *buf, T v, cuda::std::size_t n) { const auto init = blockIdx.x * blockDim.x + threadIdx.x; const auto step = blockDim.x * gridDim.x; @@ -51,7 +51,7 @@ __global__ void fill_kernel(T *buf, T v, ::cuda::std::size_t n) opts = core.ProgramOptions(include_path=str(incl.libcudacxx)) prog = core.Program(src, code_type="c++", options=opts) if data_type is None: - data_type = "::cuda::std::int32_t" + data_type = "cuda::std::int32_t" instance_name = f"fill_kernel<{data_type}>" mod = prog.compile("cubin", name_expressions=(instance_name,)) return mod.get_kernel(instance_name) diff --git a/python/examples/exec_tag_timer.py b/python/examples/exec_tag_timer.py new file mode 100644 index 0000000..2d4e093 --- /dev/null +++ b/python/examples/exec_tag_timer.py @@ -0,0 +1,143 @@ +# Copyright 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 ctypes +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: + return core.Stream.from_handle(cs.addressof()) + + +def make_copy_kernel(): + src = r""" +#include +#include + +template +__global__ void copy_kernel(const T *in, T *out, cuda::std::size_t n) +{ + const auto init = blockIdx.x * blockDim.x + threadIdx.x; + const auto step = blockDim.x * gridDim.x; + + for (auto i = init; i < n; i += step) + { + out[i] = in[i]; + } +} +""" + incl = headers.get_include_paths() + opts = core.ProgramOptions(include_path=str(incl.libcudacxx)) + prog = core.Program(src, code_type="c++", options=opts) + instance_name = "copy_kernel" + mod = prog.compile("cubin", name_expressions=(instance_name,)) + return mod.get_kernel(instance_name) + + +def make_sequence_kernel(): + src = r""" +#include +#include + +template +__global__ void sequence_kernel(T *buf, cuda::std::size_t n) +{ + const auto init = blockIdx.x * blockDim.x + threadIdx.x; + const auto step = blockDim.x * gridDim.x; + + for (auto i = init; i < n; i += step) + { + buf[i] = static_cast(i); + } +} +""" + incl = headers.get_include_paths() + opts = core.ProgramOptions(include_path=str(incl.libcudacxx)) + prog = core.Program(src, code_type="c++", options=opts) + instance_name = "sequence_kernel" + mod = prog.compile("cubin", name_expressions=(instance_name,)) + return mod.get_kernel(instance_name) + + +def make_mod2_inplace_kernel(): + src = r""" +#include +#include + +__global__ void mod2_inplace_kernel(cuda::std::int32_t *data, + cuda::std::size_t n) +{ + const auto init = blockIdx.x * blockDim.x + threadIdx.x; + const auto step = blockDim.x * gridDim.x; + + for (auto i = init; i < n; i += step) + { + data[i] = data[i] % 2; + } +} +""" + 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=("mod2_inplace_kernel",)) + return mod.get_kernel("mod2_inplace_kernel") + + +@bench.register() +def mod2_inplace(state: bench.State) -> None: + num_values = 64 * 1024 * 1024 // ctypes.sizeof(ctypes.c_int32(0)) + nbytes = num_values * ctypes.sizeof(ctypes.c_int32(0)) + + alloc_stream = as_core_Stream(state.get_stream()) + mem = core.DeviceMemoryResource(state.get_device()) + input_buf = mem.allocate(nbytes, alloc_stream) + data_buf = mem.allocate(nbytes, alloc_stream) + + state.add_element_count(num_values) + state.add_global_memory_reads(nbytes) + state.add_global_memory_writes(nbytes) + + sequence_kernel = make_sequence_kernel() + copy_kernel = make_copy_kernel() + mod2_kernel = make_mod2_inplace_kernel() + + threads_per_block = 256 + blocks_in_grid = (num_values + threads_per_block - 1) // threads_per_block + launch_config = core.LaunchConfig( + grid=blocks_in_grid, block=threads_per_block, shmem_size=0 + ) + + core.launch(alloc_stream, launch_config, sequence_kernel, input_buf, num_values) + + def launcher(launch: bench.Launch, timer: bench.Timer): + stream = as_core_Stream(launch.get_stream()) + + # Reset working data before timing the in-place operation. + core.launch(stream, launch_config, copy_kernel, input_buf, data_buf, num_values) + + timer.start() + core.launch(stream, launch_config, mod2_kernel, data_buf, num_values) + timer.stop() + + state.exec(launcher, timer=True) + + +if __name__ == "__main__": + bench.run_all_benchmarks(sys.argv) diff --git a/python/examples/skip.py b/python/examples/skip.py index 25d01d1..bfd6696 100644 --- a/python/examples/skip.py +++ b/python/examples/skip.py @@ -34,7 +34,7 @@ def make_sleep_kernel(): // Each launched thread just sleeps for `seconds`. __global__ void sleep_kernel(double seconds) { - namespace chrono = ::cuda::std::chrono; + namespace chrono = cuda::std::chrono; using hr_clock = chrono::high_resolution_clock; auto duration = static_cast(seconds * 1e9); diff --git a/python/src/py_nvbench.cpp b/python/src/py_nvbench.cpp index c872f62..10e0a0b 100644 --- a/python/src/py_nvbench.cpp +++ b/python/src/py_nvbench.cpp @@ -117,6 +117,75 @@ private: std::shared_ptr m_fn; }; +struct py_timer +{ + using callback_t = void (*)(void *); + + py_timer(void *timer, callback_t start, callback_t stop) + : m_timer{timer} + , m_start{start} + , m_stop{stop} + , m_valid{true} + {} + + void start() + { + this->check_valid(); + m_start(m_timer); + } + + void stop() + { + this->check_valid(); + m_stop(m_timer); + } + + void invalidate() noexcept + { + m_valid = false; + m_timer = nullptr; + } + +private: + void check_valid() const + { + if (!m_valid || !m_timer) + { + throw std::runtime_error("Timer is no longer valid."); + } + } + + void *m_timer{}; + callback_t m_start{}; + callback_t m_stop{}; + bool m_valid{false}; +}; + +template +py_timer make_py_timer(TimerT &timer) +{ + return py_timer{std::addressof(timer), + [](void *timer_ptr) { static_cast(timer_ptr)->start(); }, + [](void *timer_ptr) { static_cast(timer_ptr)->stop(); }}; +} + +struct py_timer_invalidation_guard +{ + explicit py_timer_invalidation_guard(py_timer &timer) + : m_timer{timer} + {} + + py_timer_invalidation_guard(const py_timer_invalidation_guard &) = delete; + py_timer_invalidation_guard(py_timer_invalidation_guard &&) = delete; + py_timer_invalidation_guard &operator=(const py_timer_invalidation_guard &) = delete; + py_timer_invalidation_guard &operator=(py_timer_invalidation_guard &&) = delete; + + ~py_timer_invalidation_guard() noexcept { m_timer.invalidate(); } + +private: + py_timer &m_timer; +}; + // Use struct to ensure public inheritance struct nvbench_run_error : std::runtime_error { @@ -340,6 +409,48 @@ void def_class_Launch(py::module_ m) py::return_value_policy::reference); } +void def_class_Timer(py::module_ m) +{ + static constexpr const char *class_Timer_doc = R"XXXX( + Controls the manually timed region of a benchmark launch. + + Each call to start() must be paired with a corresponding call to stop() + before the launch callable returns. NVBench does not validate all possible + unpaired or misordered start()/stop() sequences; benchmark results from + such use should not be trusted. + + A launch callable may call start() and stop() more than once, matching the + C++ API behavior. Repeated pairs overwrite the timer state for the launch; + they do not accumulate elapsed time and do not create additional samples. + + Note + ---- + The class is not user-constructible. NVBench provides Timer instances + to launch callables that request manual timing. +)XXXX"; + auto py_timer_cls = py::class_(m, "Timer", class_Timer_doc); + + static constexpr const char *method_start_doc = R"XXXX( + Start the timed region. + + This call must be paired with a corresponding stop() call before the launch + callable returns. Calling start()/stop() repeatedly in the same launch + overwrites the recorded interval rather than accumulating time or creating + additional samples. +)XXXX"; + py_timer_cls.def("start", &py_timer::start, method_start_doc); + + static constexpr const char *method_stop_doc = R"XXXX( + Stop the timed region. + + This records the interval since the most recent start() call. It must be + paired with a preceding start() call. Calling start()/stop() repeatedly in + the same launch overwrites the recorded interval rather than accumulating + time or creating additional samples. +)XXXX"; + py_timer_cls.def("stop", &py_timer::stop, method_stop_doc); +} + static void def_class_Benchmark(py::module_ m) { // Define Benchmark class @@ -971,13 +1082,18 @@ Use argument True to disable use of blocking kernel by NVBench" py::arg("duration_seconds")); // method State.exec - auto method_exec_impl = - [](nvbench::state &state, py::object py_launcher_fn, bool batched, bool sync) -> void { + auto method_exec_impl = [](nvbench::state &state, + py::object py_launcher_fn, + py::object py_batched, + bool sync, + bool timer) -> void { if (!PyCallable_Check(py_launcher_fn.ptr())) { throw py::type_error("Argument of exec method must be a callable object"); } + const bool batched = py_batched.is_none() ? !timer : py_batched.cast(); + // wrapper to invoke Python callable auto cpp_launcher_fn = [py_launcher_fn](nvbench::launch &launch_descr) -> void { // cast C++ object to python object @@ -986,6 +1102,36 @@ Use argument True to disable use of blocking kernel by NVBench" py_launcher_fn(launch_pyarg); }; + auto cpp_timer_launcher_fn = [py_launcher_fn](nvbench::launch &launch_descr, + auto &timer_descr) -> void { + auto launch_pyarg = py::cast(std::ref(launch_descr), py::return_value_policy::reference); + auto timer_pyarg = py::cast(make_py_timer(timer_descr)); + auto &timer_ref = timer_pyarg.template cast(); + py_timer_invalidation_guard guard{timer_ref}; + + py_launcher_fn(launch_pyarg, timer_pyarg); + }; + + if (timer) + { + if (batched) + { + throw py::value_error("State.exec(..., timer=True) requires batched=False."); + } + + if (sync) + { + constexpr auto tag = nvbench::exec_tag::timer | nvbench::exec_tag::sync; + state.exec(tag, cpp_timer_launcher_fn); + } + else + { + constexpr auto tag = nvbench::exec_tag::timer; + state.exec(tag, cpp_timer_launcher_fn); + } + return; + } + if (sync) { if (batched) @@ -1017,19 +1163,29 @@ Use argument True to disable use of blocking kernel by NVBench" Execute callable running the benchmark. The callable may be executed multiple times. The callable - will be passed `Launch` object argument. + will be passed a `Launch` object argument by default. When `timer=True`, + the callable will be passed `Launch` and `Timer` arguments. Parameters ---------- fn: Callable Python callable with signature fn(Launch) -> None that executes the benchmark. - batched: bool, optional + batched: bool or None, optional If `True`, no cache flushing is performed between callable invocations. - Default: `True`. + If `None`, defaults to `True` unless timer=True is set. When + timer=True is set, batched defaults to `False`. + Default: `None`. sync: bool, optional True value indicates that callable performs device synchronization. NVBench disables use of blocking kernel in this case. Default: `False`. + timer: bool, optional + True value requests manual timing. The callable must have signature + fn(Launch, Timer) -> None and call Timer.start() / Timer.stop() to + delimit the timed region. Passing timer=True and batched=True is + invalid because nvbench::exec_tag::timer in the C++ API disables + batched measurement. + Default: `False`. )XXXX"; pystate_cls.def("exec", @@ -1037,8 +1193,10 @@ Use argument True to disable use of blocking kernel by NVBench" method_exec_doc, py::arg("launcher_fn"), py::pos_only{}, - py::arg("batched") = true, - py::arg("sync") = false); + py::kw_only{}, + py::arg("batched") = py::none(), + py::arg("sync") = false, + py::arg("timer") = false); // method State.get_short_description static constexpr const char *method_get_short_description_doc = R"XXXX( @@ -1140,6 +1298,8 @@ PYBIND11_MODULE(PYBIND11_MODULE_NAME, m) def_class_Launch(m); + def_class_Timer(m); + def_class_Benchmark(m); def_class_State(m); diff --git a/python/test/test_cuda_bench.py b/python/test/test_cuda_bench.py index c856a5a..068a91d 100644 --- a/python/test/test_cuda_bench.py +++ b/python/test/test_cuda_bench.py @@ -32,7 +32,7 @@ def test_py_exception(): @pytest.mark.parametrize( - "cls", [bench.CudaStream, bench.State, bench.Launch, bench.Benchmark] + "cls", [bench.CudaStream, bench.State, bench.Launch, bench.Timer, bench.Benchmark] ) def test_api_ctor(cls): with pytest.raises(TypeError, match="No constructor defined!"): @@ -50,11 +50,35 @@ def t_bench(state: bench.State): def test_cpu_only(): + saved_timers = [] + + def t_bench_timer(state: bench.State): + s = {"a": 1, "b": 0.5, "c": "test", "d": {"a": 1}} + + def launcher(launch: bench.Launch, timer: bench.Timer): + saved_timers.append(timer) + timer.start() + for _ in range(10000): + _ = json.dumps(s) + timer.stop() + + with pytest.raises(ValueError, match=r"timer=True.*batched=False"): + state.exec(launcher, timer=True, batched=True) + + state.exec(launcher, timer=True) + b = bench.register(t_bench) b.set_is_cpu_only(True) + b_timer = bench.register(t_bench_timer) + b_timer.set_is_cpu_only(True) + bench.run_all_benchmarks(["-q", "--profile"]) + assert saved_timers + with pytest.raises(RuntimeError, match="Timer is no longer valid"): + saved_timers[0].start() + def docstring_check(doc_str: Union[str, None]) -> None: assert isinstance(doc_str, str) @@ -226,6 +250,13 @@ def test_Launch_doc(): obj_has_docstring_check(cl.get_stream) +def test_Timer_doc(): + cl = bench.Timer + obj_has_docstring_check(cl) + obj_has_docstring_check(cl.start) + obj_has_docstring_check(cl.stop) + + def test_CudaStream_doc(): cl = bench.CudaStream obj_has_docstring_check(cl)