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