From 281a08a57e8c79966f36f525a07630930529f3ac Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk <21087696+oleksandr-pavlyk@users.noreply.github.com> Date: Mon, 28 Jul 2025 07:51:50 -0500 Subject: [PATCH 1/7] Remove CLI --run-once and --disable-blocking-kernel options Removed option_parser::disable_blocking_kernel and option_parse::set_run_once methods. Added option_parser::enable_profile method instead, which calls ``` bench.set_run_once(true); bench.disable_blocking_kernel(true); ``` --- nvbench/option_parser.cu | 35 +++++------------------------------ nvbench/option_parser.cuh | 4 ++-- 2 files changed, 7 insertions(+), 32 deletions(-) diff --git a/nvbench/option_parser.cu b/nvbench/option_parser.cu index 4bffa0f..dce9966 100644 --- a/nvbench/option_parser.cu +++ b/nvbench/option_parser.cu @@ -456,26 +456,15 @@ void option_parser::parse_range(option_parser::arg_iterator_t first, this->lock_gpu_clocks(first[1]); first += 2; } - else if (arg == "--run-once") - { - this->enable_run_once(); - first += 1; - } else if (arg == "--stopping-criterion") { check_params(1); this->set_stopping_criterion(first[1]); first += 2; } - else if (arg == "--disable-blocking-kernel") - { - this->disable_blocking_kernel(); - first += 1; - } else if (arg == "--profile") { - this->enable_run_once(); - this->disable_blocking_kernel(); + this->enable_profile(); first += 1; } else if (arg == "--quiet" || arg == "-q") @@ -738,19 +727,6 @@ catch (std::exception &e) e.what()); } -void option_parser::enable_run_once() -{ - // If no active benchmark, save args as global. - if (m_benchmarks.empty()) - { - m_global_benchmark_args.push_back("--run-once"); - return; - } - - benchmark_base &bench = *m_benchmarks.back(); - bench.set_run_once(true); -} - void option_parser::set_stopping_criterion(const std::string &criterion) try { @@ -773,17 +749,16 @@ catch (std::exception &e) e.what()); } -void option_parser::disable_blocking_kernel() +void option_parser::enable_profile() { - // If no active benchmark, save args as global. - if (m_benchmarks.empty()) + // If no active benchmakr, save args as global { - m_global_benchmark_args.push_back("--disable-blocking-kernel"); + m_global_benchmark_args.push_back("--profile"); return; } - benchmark_base &bench = *m_benchmarks.back(); bench.set_disable_blocking_kernel(true); + bench.set_run_once(true); } void option_parser::add_benchmark(const std::string &name) diff --git a/nvbench/option_parser.cuh b/nvbench/option_parser.cuh index bb12513..110a844 100644 --- a/nvbench/option_parser.cuh +++ b/nvbench/option_parser.cuh @@ -88,8 +88,8 @@ private: void lock_gpu_clocks(const std::string &rate); void set_stopping_criterion(const std::string &criterion); - void enable_run_once(); - void disable_blocking_kernel(); + + void enable_profile(); void add_benchmark(const std::string &name); void replay_global_args(); From 3bb34b1b1f4c54791bbec32cdda9b8b7e2f76dd0 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk <21087696+oleksandr-pavlyk@users.noreply.github.com> Date: Mon, 28 Jul 2025 07:54:16 -0500 Subject: [PATCH 2/7] Remove suggestion to use --disable-blocking-kernel The text printed when blocking kernel times out already suggests to use --profile option. --- nvbench/blocking_kernel.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/nvbench/blocking_kernel.cu b/nvbench/blocking_kernel.cu index 9514ee6..d9cb975 100644 --- a/nvbench/blocking_kernel.cu +++ b/nvbench/blocking_kernel.cu @@ -83,7 +83,7 @@ __global__ void block_stream(const volatile nvbench::int32_t *flag, "NVBench documentation.\n" "\n" "If this happens while profiling with an external tool,\n" - "pass the `--disable-blocking-kernel` flag or the `--profile` flag\n" + "the `--profile` flag\n" "(to also only run the benchmark once) to the executable.\n" "\n" "For more information, see the 'Benchmark Properties' section of the\n" From 8416342af01aa91647945955b705b4554395c954 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk <21087696+oleksandr-pavlyk@users.noreply.github.com> Date: Mon, 28 Jul 2025 07:55:25 -0500 Subject: [PATCH 3/7] Remove mentions of --run-once and --disable-blocking-kernel from help Text for --profile modified to be self-consistent, i.e., not to refer to removed --run-once and --disable-blocking-kernel for explanantion of what it does. --- docs/cli_help.md | 15 ++------------- 1 file changed, 2 insertions(+), 13 deletions(-) diff --git a/docs/cli_help.md b/docs/cli_help.md index 1479ebd..d755096 100644 --- a/docs/cli_help.md +++ b/docs/cli_help.md @@ -108,25 +108,14 @@ * Applies to the most recent `--benchmark`, or all benchmarks if specified before any `--benchmark` arguments. -* `--run-once` +* `--profile` + * Don't use the `blocking_kernel`. * Only run the benchmark once, skipping any warmup runs and batched measurements. * Intended for use with external profiling tools. * Applies to the most recent `--benchmark`, or all benchmarks if specified before any `--benchmark` arguments. -* `--disable-blocking-kernel` - * Don't use the `blocking_kernel`. - * Intended for use with external profiling tools. - * Applies to the most recent `--benchmark`, or all benchmarks if specified - before any `--benchmark` arguments. - -* `--profile` - * Implies `--run-once` and `--disable-blocking-kernel`. - * Intended for use with external profiling tools. - * Applies to the most recent `--benchmark`, or all benchmarks if specified - before any `--benchmark` arguments. - ## Stopping Criteria * `--timeout ` From d160a2bafa9df627fe9d003cd1945eb4fde7c955 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk <21087696+oleksandr-pavlyk@users.noreply.github.com> Date: Mon, 28 Jul 2025 12:03:42 -0500 Subject: [PATCH 4/7] Replace --run-once in testing/CMakeLists.txt with --profile --- testing/CMakeLists.txt | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/testing/CMakeLists.txt b/testing/CMakeLists.txt index 5ee7824..8792408 100644 --- a/testing/CMakeLists.txt +++ b/testing/CMakeLists.txt @@ -31,8 +31,8 @@ set(test_srcs # Custom arguments: # CTest commands+args can't be modified after creation, so we need to rely on substitution. -set(NVBench_TEST_ARGS_nvbench.test.custom_main_custom_args "--quiet" "--my-custom-arg" "--run-once" "-d" "0") -set(NVBench_TEST_ARGS_nvbench.test.custom_main_custom_exceptions "--quiet" "--run-once" "-d" "0") +set(NVBench_TEST_ARGS_nvbench.test.custom_main_custom_args "--quiet" "--my-custom-arg" "--profile" "-d" "0") +set(NVBench_TEST_ARGS_nvbench.test.custom_main_custom_exceptions "--quiet" "--profile" "-d" "0") # Metatarget for all tests: add_custom_target(nvbench.test.all) From 5b6c3818f4ca6265260c7be33aaba9dc8f45dd8d Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk <21087696+oleksandr-pavlyk@users.noreply.github.com> Date: Mon, 28 Jul 2025 14:39:54 -0500 Subject: [PATCH 5/7] Corrected blocking kernel timeout message --- nvbench/blocking_kernel.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/nvbench/blocking_kernel.cu b/nvbench/blocking_kernel.cu index d9cb975..f9a24a1 100644 --- a/nvbench/blocking_kernel.cu +++ b/nvbench/blocking_kernel.cu @@ -83,8 +83,8 @@ __global__ void block_stream(const volatile nvbench::int32_t *flag, "NVBench documentation.\n" "\n" "If this happens while profiling with an external tool,\n" - "the `--profile` flag\n" - "(to also only run the benchmark once) to the executable.\n" + "pass the `--profile` flag to the executable to disable use of blocking kernel\n" + "and to also run the benchmark only once.\n" "\n" "For more information, see the 'Benchmark Properties' section of the\n" "NVBench documentation.\n\n", From 25c604cf37565b88b8f6e0f2fe5369b78eba94be Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk <21087696+oleksandr-pavlyk@users.noreply.github.com> Date: Mon, 28 Jul 2025 14:40:02 -0500 Subject: [PATCH 6/7] Fix typo, corrected control flow logic flaw Verify that `--profile` option results in setting `m_run_once`: ``` (nvbench) opavlyk@ee09c48-lcedt:~/repos/nvbench$ ./build/bin/nvbench.example.cpp20.axes -b simple -d 0 --profile | grep "Pass: Cold" Pass: Cold: 1.006560ms GPU, 1.009277ms CPU, 0.00s total GPU, 0.00s total wall, 1x (nvbench) opavlyk@ee09c48-lcedt:~/repos/nvbench$ ./build/bin/nvbench.example.cpp20.axes -b simple -d 0 | grep "Pass: Cold" Pass: Cold: 1.002844ms GPU, 1.011917ms CPU, 0.50s total GPU, 0.52s total wall, 499x ``` --- nvbench/option_parser.cu | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/nvbench/option_parser.cu b/nvbench/option_parser.cu index dce9966..efddbb4 100644 --- a/nvbench/option_parser.cu +++ b/nvbench/option_parser.cu @@ -751,7 +751,8 @@ catch (std::exception &e) void option_parser::enable_profile() { - // If no active benchmakr, save args as global + // If no active benchmark, save args as global + if (m_benchmarks.empty()) { m_global_benchmark_args.push_back("--profile"); return; From 4ad3088a47195ea0a858a8e2fecfbf7b6665bed2 Mon Sep 17 00:00:00 2001 From: Oleksandr Pavlyk <21087696+oleksandr-pavlyk@users.noreply.github.com> Date: Mon, 28 Jul 2025 14:52:57 -0500 Subject: [PATCH 7/7] Update docs/cli_help.md Spare users of implementation details in description of `--profile` option Co-authored-by: Allison Piper --- docs/cli_help.md | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/docs/cli_help.md b/docs/cli_help.md index d755096..35265a1 100644 --- a/docs/cli_help.md +++ b/docs/cli_help.md @@ -109,9 +109,8 @@ before any `--benchmark` arguments. * `--profile` - * Don't use the `blocking_kernel`. - * Only run the benchmark once, skipping any warmup runs and batched - measurements. + * Only run each benchmark once. + * Disable any instrumentation that may interfere with profilers. * Intended for use with external profiling tools. * Applies to the most recent `--benchmark`, or all benchmarks if specified before any `--benchmark` arguments.