From 165cf924c5b04d4ff8a5e8de08e35928b5e64c4b Mon Sep 17 00:00:00 2001 From: Allison Piper Date: Tue, 9 Apr 2024 12:45:58 -0400 Subject: [PATCH] Refactor main implementation to improve reusability and customization. (#165) * Refactor main implementation to improve reusability and customization. Move the implementation of `main` out of macros and into separate functions. This allows for easier reuse and customization of the macros. Existing macro usage should still work as expected, and new customization points will simplify common tasks like argument parsing going forward. * Add tests that validate common main customizations. --- nvbench/main.cuh | 241 ++++++++++++++++++----- testing/CMakeLists.txt | 15 +- testing/custom_main_custom_args.cu | 132 +++++++++++++ testing/custom_main_custom_exceptions.cu | 64 ++++++ 4 files changed, 403 insertions(+), 49 deletions(-) create mode 100644 testing/custom_main_custom_args.cu create mode 100644 testing/custom_main_custom_exceptions.cu diff --git a/nvbench/main.cuh b/nvbench/main.cuh index bcdced5..84b8ebb 100644 --- a/nvbench/main.cuh +++ b/nvbench/main.cuh @@ -28,14 +28,130 @@ #include #include +// Advanced users can rebuild NVBench's `main` function using the macros in this file, or replace +// them with customized implementations. + +// Customization point, called before NVBench initialization. +#ifndef NVBENCH_MAIN_INITIALIZE_CUSTOM_PRE +#define NVBENCH_MAIN_INITIALIZE_CUSTOM_PRE(argc, argv) []() {}() +#endif + +// Customization point, called after NVBench initialization. +#ifndef NVBENCH_MAIN_INITIALIZE_CUSTOM_POST +#define NVBENCH_MAIN_INITIALIZE_CUSTOM_POST(argc, argv) []() {}() +#endif + +// Customization point, called before NVBench parsing. Update argc/argv if needed. +// argc/argv are the usual command line arguments types. The ARGS version of this +// macro is a bit more convenient. +#ifndef NVBENCH_MAIN_CUSTOM_ARGC_ARGV_HANDLER +#define NVBENCH_MAIN_CUSTOM_ARGC_ARGV_HANDLER(argc, argv) []() {}() +#endif + +// Customization point, called before NVBench parsing. Update args if needed. +// Args is a vector of strings, each element is an argument. +#ifndef NVBENCH_MAIN_CUSTOM_ARGS_HANDLER +#define NVBENCH_MAIN_CUSTOM_ARGS_HANDLER(args) []() {}() +#endif + +// Customization point, called before NVBench parsing. +#ifndef NVBENCH_MAIN_PARSE_CUSTOM_PRE +#define NVBENCH_MAIN_PARSE_CUSTOM_PRE(parser, args) []() {}() +#endif + +// Customization point, called after NVBench parsing. +#ifndef NVBENCH_MAIN_PARSE_CUSTOM_POST +#define NVBENCH_MAIN_PARSE_CUSTOM_POST(parser) []() {}() +#endif + +// Customization point, called before NVBench finalization. +#ifndef NVBENCH_MAIN_FINALIZE_CUSTOM_PRE +#define NVBENCH_MAIN_FINALIZE_CUSTOM_PRE() []() {}() +#endif + +// Customization point, called after NVBench finalization. +#ifndef NVBENCH_MAIN_FINALIZE_CUSTOM_POST +#define NVBENCH_MAIN_FINALIZE_CUSTOM_POST() []() {}() +#endif + +// Customization point, use to catch addition exceptions. +#ifndef NVBENCH_MAIN_CATCH_EXCEPTIONS_CUSTOM +#define NVBENCH_MAIN_CATCH_EXCEPTIONS_CUSTOM +#endif + +/************************************ Default implementation **************************************/ + +#ifndef NVBENCH_MAIN #define NVBENCH_MAIN \ - int main(int argc, char const *const *argv) \ + int main(int argc, char **argv) \ try \ { \ NVBENCH_MAIN_BODY(argc, argv); \ - NVBENCH_CUDA_CALL(cudaDeviceReset()); \ return 0; \ } \ + NVBENCH_MAIN_CATCH_EXCEPTIONS_CUSTOM \ + NVBENCH_MAIN_CATCH_EXCEPTIONS +#endif + +#ifndef NVBENCH_MAIN_BODY +#define NVBENCH_MAIN_BODY(argc, argv) \ + NVBENCH_MAIN_INITIALIZE(argc, argv); \ + { \ + NVBENCH_MAIN_PARSE(argc, argv); \ + \ + NVBENCH_MAIN_PRINT_PREAMBLE(parser); \ + NVBENCH_MAIN_RUN_BENCHMARKS(parser); \ + NVBENCH_MAIN_PRINT_EPILOGUE(parser); \ + \ + NVBENCH_MAIN_PRINT_RESULTS(parser); \ + } /* Tear down parser before finalization */ \ + NVBENCH_MAIN_FINALIZE(); \ + return 0; +#endif + +#ifndef NVBENCH_MAIN_INITIALIZE +#define NVBENCH_MAIN_INITIALIZE(argc, argv) \ + NVBENCH_MAIN_INITIALIZE_CUSTOM_PRE(argc, argv); \ + nvbench::detail::main_initialize(argc, argv); \ + NVBENCH_MAIN_INITIALIZE_CUSTOM_POST(argc, argv) +#endif + +#ifndef NVBENCH_MAIN_PARSE +#define NVBENCH_MAIN_PARSE(argc, argv) \ + NVBENCH_MAIN_CUSTOM_ARGC_ARGV_HANDLER(argc, argv); \ + std::vector args = nvbench::detail::main_convert_args(argc, argv); \ + NVBENCH_MAIN_CUSTOM_ARGS_HANDLER(args); \ + nvbench::option_parser parser; \ + NVBENCH_MAIN_PARSE_CUSTOM_PRE(parser, args); \ + parser.parse(args); \ + NVBENCH_MAIN_PARSE_CUSTOM_POST(parser) +#endif + +#ifndef NVBENCH_MAIN_PRINT_PREAMBLE +#define NVBENCH_MAIN_PRINT_PREAMBLE(parser) nvbench::detail::main_print_preamble(parser) +#endif + +#ifndef NVBENCH_MAIN_RUN_BENCHMARKS +#define NVBENCH_MAIN_RUN_BENCHMARKS(parser) nvbench::detail::main_run_benchmarks(parser) +#endif + +#ifndef NVBENCH_MAIN_PRINT_EPILOGUE +#define NVBENCH_MAIN_PRINT_EPILOGUE(parser) nvbench::detail::main_print_epilogue(parser) +#endif + +#ifndef NVBENCH_MAIN_PRINT_RESULTS +#define NVBENCH_MAIN_PRINT_RESULTS(parser) nvbench::detail::main_print_results(parser) +#endif + +#ifndef NVBENCH_MAIN_FINALIZE +#define NVBENCH_MAIN_FINALIZE() \ + NVBENCH_MAIN_FINALIZE_CUSTOM_PRE(); \ + nvbench::detail::main_finalize(); \ + NVBENCH_MAIN_FINALIZE_CUSTOM_POST() +#endif + +#ifndef NVBENCH_MAIN_CATCH_EXCEPTIONS +#define NVBENCH_MAIN_CATCH_EXCEPTIONS \ catch (std::exception & e) \ { \ std::cerr << "\nNVBench encountered an error:\n\n" << e.what() << "\n"; \ @@ -46,56 +162,87 @@ std::cerr << "\nNVBench encountered an unknown error.\n"; \ return 1; \ } - -#ifdef NVBENCH_HAS_CUPTI -#define NVBENCH_INITIALIZE_DRIVER_API NVBENCH_DRIVER_API_CALL(cuInit(0)) -#else -// clang-format off -#define NVBENCH_INITIALIZE_DRIVER_API do {} while (false) -// clang-format on #endif -#define NVBENCH_MAIN_PARSE(argc, argv) \ - nvbench::option_parser parser; \ - parser.parse(argc, argv) +namespace nvbench::detail +{ -// See NVIDIA/NVBench#136 for CUDA_MODULE_LOADING +inline void set_env(const char *name, const char *value) +{ #ifdef _MSC_VER -#define NVBENCH_INITIALIZE_CUDA_ENV _putenv_s("CUDA_MODULE_LOADING", "EAGER") + _putenv_s(name, value); #else -#define NVBENCH_INITIALIZE_CUDA_ENV setenv("CUDA_MODULE_LOADING", "EAGER", 1) + setenv(name, value, 1); +#endif +} + +inline void main_initialize(int, char **) +{ + // See NVIDIA/NVBench#136 for CUDA_MODULE_LOADING + set_env("CUDA_MODULE_LOADING", "EAGER"); + + // Initialize CUDA driver API if needed: +#ifdef NVBENCH_HAS_CUPTI + NVBENCH_DRIVER_API_CALL(cuInit(0)); #endif -#define NVBENCH_INITIALIZE_BENCHMARKS() \ - nvbench::benchmark_manager::get().initialize() + // Initialize the benchmarks *after* setting up the CUDA environment: + nvbench::benchmark_manager::get().initialize(); +} -#define NVBENCH_MAIN_BODY(argc, argv) \ - do \ - { \ - NVBENCH_INITIALIZE_CUDA_ENV; \ - NVBENCH_INITIALIZE_DRIVER_API; \ - NVBENCH_INITIALIZE_BENCHMARKS(); \ - NVBENCH_MAIN_PARSE(argc, argv); \ - auto &printer = parser.get_printer(); \ - \ - printer.print_device_info(); \ - printer.print_log_preamble(); \ - auto &benchmarks = parser.get_benchmarks(); \ - \ - std::size_t total_states = 0; \ - for (auto &bench_ptr : benchmarks) \ - { \ - total_states += bench_ptr->get_config_count(); \ - } \ - printer.set_total_state_count(total_states); \ - \ - printer.set_completed_state_count(0); \ - for (auto &bench_ptr : benchmarks) \ - { \ - bench_ptr->set_printer(printer); \ - bench_ptr->run(); \ - bench_ptr->clear_printer(); \ - } \ - printer.print_log_epilogue(); \ - printer.print_benchmark_results(benchmarks); \ - } while (false) +inline std::vector main_convert_args(int argc, char const *const *argv) +{ + std::vector args; + for (int i = 0; i < argc; ++i) + { + args.push_back(argv[i]); + } + return args; +} + +inline void main_print_preamble(option_parser &parser) +{ + auto &printer = parser.get_printer(); + + printer.print_device_info(); + printer.print_log_preamble(); +} + +inline void main_run_benchmarks(option_parser &parser) +{ + auto &printer = parser.get_printer(); + auto &benchmarks = parser.get_benchmarks(); + + std::size_t total_states = 0; + for (auto &bench_ptr : benchmarks) + { + total_states += bench_ptr->get_config_count(); + } + + printer.set_completed_state_count(0); + printer.set_total_state_count(total_states); + + for (auto &bench_ptr : benchmarks) + { + bench_ptr->set_printer(printer); + bench_ptr->run(); + bench_ptr->clear_printer(); + } +} + +inline void main_print_epilogue(option_parser &parser) +{ + auto &printer = parser.get_printer(); + printer.print_log_epilogue(); +} + +inline void main_print_results(option_parser &parser) +{ + auto &printer = parser.get_printer(); + auto &benchmarks = parser.get_benchmarks(); + printer.print_benchmark_results(benchmarks); +} + +inline void main_finalize() { NVBENCH_CUDA_CALL(cudaDeviceReset()); } + +} // namespace nvbench::detail diff --git a/testing/CMakeLists.txt b/testing/CMakeLists.txt index 1535878..55eb741 100644 --- a/testing/CMakeLists.txt +++ b/testing/CMakeLists.txt @@ -6,6 +6,8 @@ set(test_srcs cpu_timer.cu criterion_manager.cu criterion_params.cu + custom_main_custom_args.cu + custom_main_custom_exceptions.cu enum_type_list.cu entropy_criterion.cu float64_axis.cu @@ -24,7 +26,12 @@ set(test_srcs type_list.cu ) -# Metatarget for all examples: +# 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") + +# Metatarget for all tests: add_custom_target(nvbench.test.all) add_dependencies(nvbench.all nvbench.test.all) @@ -36,10 +43,14 @@ foreach(test_src IN LISTS test_srcs) target_link_libraries(${test_name} PRIVATE nvbench::nvbench fmt) set_target_properties(${test_name} PROPERTIES COMPILE_FEATURES cuda_std_17) nvbench_config_target(${test_name}) - add_test(NAME ${test_name} COMMAND "$") + add_test(NAME ${test_name} COMMAND "$" ${NVBench_TEST_ARGS_${test_name}}) add_dependencies(nvbench.test.all ${test_name}) endforeach() +set_tests_properties(nvbench.test.custom_main_custom_exceptions PROPERTIES + PASS_REGULAR_EXPRESSION "Custom error detected: Expected exception thrown." +) + add_subdirectory(cmake) add_subdirectory(device) diff --git a/testing/custom_main_custom_args.cu b/testing/custom_main_custom_args.cu new file mode 100644 index 0000000..f7e331e --- /dev/null +++ b/testing/custom_main_custom_args.cu @@ -0,0 +1,132 @@ +/* + * Copyright 2024 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. + */ + +#include +#include "nvbench/cuda_call.cuh" + +/****************************************************************************** + * Install custom parser. + * sSee for more details. + ******************************************************************************/ + +// +// Step 1: Define a custom argument handler that accepts a vector of strings. +// - This handler should modify the vector in place to remove any custom +// arguments it handles. NVbench will then parse the remaining arguments. +// - The handler should also update any application state needed to handle +// the custom arguments. +// + +// User code to handle a specific argument: +void handle_my_custom_arg(); + +// NVBench hook for modiifying the command line arguments before parsing: +void custom_arg_handler(std::vector &args) +{ + // Handle and remove "--my-custom-arg" + if (auto it = std::find(args.begin(), args.end(), "--my-custom-arg"); it != args.end()) + { + handle_my_custom_arg(); + args.erase(it); + } +} + +// +// Step 2: Install the custom argument handler. +// - This is done by defining a macro that invokes the custom argument handler. +// + +// Install the custom argument handler: +// Either define this before any NVBench headers are included, or undefine and redefine: +#undef NVBENCH_MAIN_CUSTOM_ARGS_HANDLER +#define NVBENCH_MAIN_CUSTOM_ARGS_HANDLER(args) custom_arg_handler(args) + +// Step 3: Define `main` +// +// After installing the custom argument handler, define the main function using: +// +// ``` +// NVBENCH_MAIN +// ``` +// +// Here, this is done at the end of this file. + +/****************************************************************************** + * Unit test verification: + ******************************************************************************/ + +// Track whether the args are found / handled. +bool h_custom_arg_found = false; +bool h_handled_on_device = false; +__device__ bool d_custom_arg_found = false; +__device__ bool d_handled_on_device = false; + +// Copy host values to device: +void copy_host_state_to_device() +{ + NVBENCH_CUDA_CALL(cudaMemcpyToSymbol(d_custom_arg_found, &h_custom_arg_found, sizeof(bool))); + NVBENCH_CUDA_CALL(cudaMemcpyToSymbol(d_handled_on_device, &h_handled_on_device, sizeof(bool))); +} + +// Copy device values to host: +void copy_device_state_to_host() +{ + NVBENCH_CUDA_CALL(cudaMemcpyFromSymbol(&h_custom_arg_found, d_custom_arg_found, sizeof(bool))); + NVBENCH_CUDA_CALL(cudaMemcpyFromSymbol(&h_handled_on_device, d_handled_on_device, sizeof(bool))); +} + +void handle_my_custom_arg() +{ + h_custom_arg_found = true; + copy_host_state_to_device(); +} + +void verify() +{ + copy_device_state_to_host(); + if (!h_custom_arg_found) + { + throw std::runtime_error("Custom argument not detected."); + } + if (!h_handled_on_device) + { + throw std::runtime_error("Custom argument not handled on device."); + } +} + +// Install a verification check to ensure the custom argument was handled. +// Use the `PRE` finalize hook to ensure we check device state before resetting the context. +#undef NVBENCH_MAIN_FINALIZE_CUSTOM_PRE +#define NVBENCH_MAIN_FINALIZE_CUSTOM_PRE() verify() + +// Simple kernel/benchmark to make sure that the handler can successfully modify CUDA state: +__global__ void kernel() +{ + if (d_custom_arg_found) + { + d_handled_on_device = true; + } +} +void bench(nvbench::state &state) +{ + state.exec([](nvbench::launch &) { kernel<<<1, 1>>>(); }); +} +NVBENCH_BENCH(bench); + +// Define the customized main function: +NVBENCH_MAIN diff --git a/testing/custom_main_custom_exceptions.cu b/testing/custom_main_custom_exceptions.cu new file mode 100644 index 0000000..b1f9b9c --- /dev/null +++ b/testing/custom_main_custom_exceptions.cu @@ -0,0 +1,64 @@ +/* + * Copyright 2022 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. + */ + +#include + +#include + +/****************************************************************************** + * Install exception handler around the NVBench main body. This is used + * to print helpful information when a user exception is thrown before exiting. + * + * Note that this will **NOT** be used when a benchmark throws an exception. + * That will fail the benchmark and note the exception, and continue + * execution. + * + * This is used to catch exceptions in user extensions of NVBench, things like + * customized initialization, command line parsing, finalization, etc. See + * for more details. + ******************************************************************************/ + +struct user_exception : public std::runtime_error +{ + user_exception() + : std::runtime_error("Expected exception thrown.") + {} +}; + +// User code to handle user exception: +void handle_my_exception(user_exception &e) +{ + std::cerr << "Custom error detected: " << e.what() << std::endl; + std::exit(1); +} + +// Install the exception handler around the NVBench main body. +// NVBench will have sensible defaults for common exceptions following this if no terminating catch +// block is defined. +// Either define this before any NVBench headers are included, or undefine and redefine. +#undef NVBENCH_MAIN_CATCH_EXCEPTIONS_CUSTOM +#define NVBENCH_MAIN_CATCH_EXCEPTIONS_CUSTOM \ + catch (user_exception & e) { handle_my_exception(e); } + +// For testing purposes, install a argument parser that throws: +void really_robust_argument_parser(std::vector &) { throw user_exception(); } +#undef NVBENCH_MAIN_CUSTOM_ARGS_HANDLER +#define NVBENCH_MAIN_CUSTOM_ARGS_HANDLER(args) really_robust_argument_parser(args); + +// Define the customized main function: +NVBENCH_MAIN