From 60c94d9ed65b3ea252e47c69030aa119ba9c676f Mon Sep 17 00:00:00 2001 From: Allison Vacanti Date: Tue, 16 Mar 2021 13:57:52 -0400 Subject: [PATCH] Add `enum_type_axis` and `examples/enums.cu`. - `enum_type_axis` simplifies using integral_constants with type axes. - `examples/enums.cu` demonstrates various ways of implementing parameter sweeps with enum types. --- README.md | 1 + docs/benchmarks.md | 7 ++ examples/CMakeLists.txt | 1 + examples/enums.cu | 214 +++++++++++++++++++++++++++++++++++++ nvbench/enum_type_list.cuh | 61 +++++++++++ nvbench/nvbench.cuh | 1 + nvbench/type_strings.cu | 12 ++- nvbench/type_strings.cuh | 19 ++++ testing/CMakeLists.txt | 1 + testing/enum_type_list.cu | 189 ++++++++++++++++++++++++++++++++ 10 files changed, 503 insertions(+), 3 deletions(-) create mode 100644 examples/enums.cu create mode 100644 nvbench/enum_type_list.cuh create mode 100644 testing/enum_type_list.cu diff --git a/README.md b/README.md index c4114fe..02fe877 100644 --- a/README.md +++ b/README.md @@ -57,6 +57,7 @@ This repository provides a number of [examples](examples/) that demonstrate various NVBench features and usecases: - [Runtime and compile-time parameter sweeps](examples/axes.cu) +- [Enums and compile-time-constant-integral parameter axes](examples/enums.cu) - [Reporting item/sec and byte/sec throughput statistics](examples/throughput.cu) - [Skipping benchmark configurations](examples/skip.cu) - [Benchmarks that sync CUDA devices: `nvbench::exec_tag::sync`](examples/exec_tag_sync.cu) diff --git a/docs/benchmarks.md b/docs/benchmarks.md index e097cde..09820f9 100644 --- a/docs/benchmarks.md +++ b/docs/benchmarks.md @@ -130,6 +130,9 @@ void benchmark(nvbench::state& state) NVBENCH_BENCH(benchmark).add_string_axis("RNG Distribution", {"Uniform", "Gaussian"}); ``` +A common use for string axes is to encode enum values, as shown in +[examples/enums.cu](../examples/enums.cu). + ## Type Axes Another common situation involves benchmarking a templated kernel with multiple @@ -161,6 +164,10 @@ NVBENCH_BENCH_TYPES(my_benchmark, NVBENCH_TYPE_AXES(my_types)) The `NVBENCH_TYPE_AXES` macro is unfortunately necessary to prevent commas in the `type_list<...>` from breaking macro parsing. +Type axes can be used to encode compile-time enum and integral constants using +the `nvbench::enum_type_list` helper. See +[examples/enums.cu](../examples/enums.cu) for detail. + ## `nvbench::range` Since parameter sweeps often explore a range of evenly-spaced numeric values, a diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index d46cb82..e62decb 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -1,5 +1,6 @@ set(example_srcs axes.cu + enums.cu exec_tag_sync.cu exec_tag_timer.cu skip.cu diff --git a/examples/enums.cu b/examples/enums.cu new file mode 100644 index 0000000..e4a21d6 --- /dev/null +++ b/examples/enums.cu @@ -0,0 +1,214 @@ +/* + * Copyright 2021 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 + +// Enum to use as parameter axis: +enum class MyEnum +{ + ValueA, + ValueB, + ValueC +}; + +//============================================================================== +// Sweep through enum values at runtime using a string axis. +// +// Preferred way to provide access to an enum value that doesn't need to be +// compile-time constant. Create a string axis with unique values for each +// enum value of interest, and convert the string to an enum value in the +// benchmark. +// +// This approach is preferred since it gives nicer output (readable names +// instead of integral enum values), but takes a bit of extra work to convert +// the strings back to an enum value. +// +// `--list` output: +// ``` +// * `MyEnum` : string +// * `A` +// * `B` +// * `C` +// ``` +void runtime_enum_sweep_string(nvbench::state &state) +{ + const auto enum_string = state.get_string("MyEnum"); + [[maybe_unused]] MyEnum enum_value{}; + if (enum_string == "A") + { + enum_value = MyEnum::ValueA; + } + else if (enum_string == "B") + { + enum_value = MyEnum::ValueB; + } + else if (enum_string == "C") + { + enum_value = MyEnum::ValueC; + } + + // Do stuff with enum_value. + // Create inputs, etc, configure runtime kernel parameters, etc. + + // Just a dummy kernel. + state.exec([](nvbench::launch &launch) { + nvbench::sleep_kernel<<<1, 1, 0, launch.get_stream()>>>(1e-3); + }); +} +NVBENCH_BENCH(runtime_enum_sweep_string) + .add_string_axis("MyEnum", {"A", "B", "C"}); + +//============================================================================== +// Sweep through enum values at runtime using an int64 axis. +// +// This may be useful for doing quick tests / prototyping, but does not provide +// readable output / command-line args since numeric values will be used for the +// axis. +// +// `--list` output: +// ``` +// * `MyEnum` : int64 +// * `0` +// * `1` +// * `2` +// ``` +void runtime_enum_sweep_int64(nvbench::state &state) +{ + const auto enum_value = static_cast(state.get_int64("MyEnum")); + + // Do stuff with enum_value. + // Create inputs, etc, configure runtime kernel parameters, etc. + + // Just a dummy kernel. + state.exec([](nvbench::launch &launch) { + nvbench::sleep_kernel<<<1, 1, 0, launch.get_stream()>>>(1e-3); + }); +} +NVBENCH_BENCH(runtime_enum_sweep_int64) + .add_int64_axis("MyEnum", + {static_cast(MyEnum::ValueA), + static_cast(MyEnum::ValueB), + static_cast(MyEnum::ValueC)}); + +//============================================================================== +// Sweep through enum values at compile time using an `enum_type_axis`. +// +// If an enum value needs to be available at compile time (for example, if it's +// used as a template parameter), the `nvbench::enum_type_axis` helper can be +// used to create a type axis of `std::integral_constant` types. +// +// The `NVBENCH_DECLARE_ENUM_TYPE_STRINGS(T, InputGenerator, DescGenerator)` +// utility configures an `nvbench::type_strings` specialization for the integral +// constants, improving readability of input/output, as shown below. +// +// `--list` output: +// ``` +// * `MyEnum` : type +// * `A` (MyEnum::ValueA) +// * `B` (MyEnum::ValueB) +// * `C` (MyEnum::ValueC) +// ``` + +// Tell NVBench how to turn your enum into strings for display, commandline +// args, `--list` output, tables, etc. +NVBENCH_DECLARE_ENUM_TYPE_STRINGS( + // Enum type: + MyEnum, + // Callable to generate input strings: + // Short identifier used for tables, command-line args, etc. + // Used when context is available to figure out the enum type. + [](MyEnum value) { + switch (value) + { + case MyEnum::ValueA: + return "A"; + case MyEnum::ValueB: + return "B"; + case MyEnum::ValueC: + return "C"; + default: + return "Unknown"; + } + }, + // Callable to generate descriptions: + // If non-empty, these are used in `--list` to describe values. + // Used when context may not be available to figure out the type from the + // input string. + // Just use `[](auto) { return std::string{}; }` if you don't want these. + [](MyEnum value) { + switch (value) + { + case MyEnum::ValueA: + return "MyEnum::ValueA"; + case MyEnum::ValueB: + return "MyEnum::ValueB"; + case MyEnum::ValueC: + return "MyEnum::ValueC"; + default: + return "Unknown"; + } + }) + +// The actual compile-time enum sweep benchmark: +template +void compile_time_enum_sweep( + nvbench::state &state, + nvbench::type_list>) +{ + // Use EnumValue in compile-time contexts. + // Template parameters, static dispatch, etc. + + // Just a dummy kernel. + state.exec([](nvbench::launch &launch) { + nvbench::sleep_kernel<<<1, 1, 0, launch.get_stream()>>>(1e-3); + }); +} +using MyEnumList = + nvbench::enum_type_list; +NVBENCH_BENCH_TYPES(compile_time_enum_sweep, NVBENCH_TYPE_AXES(MyEnumList)) + .set_type_axes_names({"MyEnum"}); + +//============================================================================== +// `enum_type_list` works for other integral types, too. +// +// `--list` output: +// ``` +// * `SomeInts` : type +// * `0` (struct std::integral_constant) +// * `16` (struct std::integral_constant) +// * `4096` (struct std::integral_constant) +// * `-12` (struct std::integral_constant) +// ``` +template +void compile_time_int_sweep( + nvbench::state &state, + nvbench::type_list>) +{ + // Use IntValue in compile time contexts. + // Template parameters, static dispatch, etc. + + // Just a dummy kernel. + state.exec([](nvbench::launch &launch) { + nvbench::sleep_kernel<<<1, 1, 0, launch.get_stream()>>>(1e-3); + }); +} +using MyInts = nvbench::enum_type_list; +NVBENCH_BENCH_TYPES(compile_time_int_sweep, NVBENCH_TYPE_AXES(MyInts)) + .set_type_axes_names({"SomeInts"}); diff --git a/nvbench/enum_type_list.cuh b/nvbench/enum_type_list.cuh new file mode 100644 index 0000000..1a9bef6 --- /dev/null +++ b/nvbench/enum_type_list.cuh @@ -0,0 +1,61 @@ +/* + * Copyright 2021 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. + */ + +#pragma once + +#include +#include + +#include + +namespace nvbench +{ + +/*! + * \brief Helper utility that generates a `type_list` of + * `std::integral_constant`s. + * + * \relatesalso NVBENCH_DECLARE_ENUM_TYPE_STRINGS + */ +template +using enum_type_list = nvbench::type_list...>; +} // namespace nvbench + +/*! + * \brief Declare `type_string`s for an `enum_type_list`. + * + * Given an enum type `T` and two callables that produce input and description + * strings, declare a specialization for + * `nvbench::type_string>`. + * + * Must be used from global namespace scope. + * + * \relatesalso nvbench::enum_type_list + */ +#define NVBENCH_DECLARE_ENUM_TYPE_STRINGS(T, \ + input_generator, \ + description_generator) \ + namespace nvbench \ + { \ + template \ + struct type_strings> \ + { \ + static std::string input_string() { return input_generator(Value); } \ + static std::string description() { return description_generator(Value); } \ + }; \ + } diff --git a/nvbench/nvbench.cuh b/nvbench/nvbench.cuh index 6226f49..269575a 100644 --- a/nvbench/nvbench.cuh +++ b/nvbench/nvbench.cuh @@ -27,6 +27,7 @@ #include #include #include +#include #include #include #include diff --git a/nvbench/type_strings.cu b/nvbench/type_strings.cu index 50f4b99..261ad88 100644 --- a/nvbench/type_strings.cu +++ b/nvbench/type_strings.cu @@ -18,6 +18,8 @@ #include +#include + #include #if defined(__GNUC__) || defined(__clang__) @@ -34,13 +36,15 @@ namespace { struct free_wrapper { - void operator()(void* ptr) { std::free(ptr); } + void operator()(void *ptr) { std::free(ptr); } }; } // end namespace #endif // NVBENCH_CXXABI_DEMANGLE -namespace nvbench::detail +namespace nvbench +{ +namespace detail { std::string demangle(const std::string &str) @@ -54,4 +58,6 @@ std::string demangle(const std::string &str) #endif }; -} // namespace nvbench::detail +} // namespace detail + +} // namespace nvbench diff --git a/nvbench/type_strings.cuh b/nvbench/type_strings.cuh index 9086377..522f9f0 100644 --- a/nvbench/type_strings.cuh +++ b/nvbench/type_strings.cuh @@ -21,6 +21,7 @@ #include #include +#include #include namespace nvbench @@ -49,8 +50,26 @@ struct type_strings static std::string description() { return {}; } }; +template +struct type_strings> +{ + // The string used to identify the type in shorthand (e.g. output tables and + // CLI options): + static std::string input_string() { return std::to_string(Value); } + + // A more descriptive identifier for the type, if input_string is not a common + // identifier. May be blank if `input_string` is obvious. + static std::string description() + { + return nvbench::detail::demangle>(); + } +}; + } // namespace nvbench +/*! + * Declare an `input_string` and `description` to use with a specific `type`. + */ #define NVBENCH_DECLARE_TYPE_STRINGS(Type, InputString, Description) \ namespace nvbench \ { \ diff --git a/testing/CMakeLists.txt b/testing/CMakeLists.txt index 1b4b31c..a983dba 100644 --- a/testing/CMakeLists.txt +++ b/testing/CMakeLists.txt @@ -4,6 +4,7 @@ set(test_srcs create.cu cuda_timer.cu cpu_timer.cu + enum_type_list.cu float64_axis.cu int64_axis.cu named_values.cu diff --git a/testing/enum_type_list.cu b/testing/enum_type_list.cu new file mode 100644 index 0000000..2dd99d7 --- /dev/null +++ b/testing/enum_type_list.cu @@ -0,0 +1,189 @@ +/* + * Copyright 2021 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 "test_asserts.cuh" + +#include + +#include + +enum class scoped_enum +{ + val_1, + val_2, + val_3 +}; +NVBENCH_DECLARE_ENUM_TYPE_STRINGS( + scoped_enum, + [](scoped_enum value) { + switch (value) + { + case scoped_enum::val_1: + return fmt::format("1"); + case scoped_enum::val_2: + return fmt::format("2"); + case scoped_enum::val_3: + return fmt::format("3"); + default: + return std::string{"Unknown"}; + } + }, + [](scoped_enum value) { + switch (value) + { + case scoped_enum::val_1: + return fmt::format("scoped_enum::val_1"); + case scoped_enum::val_2: + return fmt::format("scoped_enum::val_2"); + case scoped_enum::val_3: + return fmt::format("scoped_enum::val_3"); + default: + return std::string{"Unknown"}; + } + }) + +enum unscoped_enum +{ + unscoped_val_1, + unscoped_val_2, + unscoped_val_3 +}; +NVBENCH_DECLARE_ENUM_TYPE_STRINGS( + unscoped_enum, + [](unscoped_enum value) { + switch (value) + { + case unscoped_val_1: + return fmt::format("1"); + case unscoped_val_2: + return fmt::format("2"); + case unscoped_val_3: + return fmt::format("3"); + default: + return std::string{"Unknown"}; + } + }, + [](unscoped_enum value) { + switch (value) + { + case unscoped_val_1: + return fmt::format("unscoped_val_1"); + case unscoped_val_2: + return fmt::format("unscoped_val_2"); + case unscoped_val_3: + return fmt::format("unscoped_val_3"); + default: + return std::string{"Unknown"}; + } + }) + +void test_int() +{ + ASSERT((std::is_same_v, nvbench::type_list<>>)); + ASSERT((std::is_same_v, + nvbench::type_list>>)); + ASSERT((std::is_same_v, + nvbench::type_list, + std::integral_constant, + std::integral_constant, + std::integral_constant, + std::integral_constant>>)); +} + +void test_scoped_enum() +{ + ASSERT(( + std::is_same_v, nvbench::type_list<>>)); + ASSERT(( + std::is_same_v, + nvbench::type_list< + std::integral_constant>>)); + ASSERT(( + std::is_same_v, + nvbench::type_list< + std::integral_constant, + std::integral_constant, + std::integral_constant>>)); +} + +void test_unscoped_enum() +{ + ASSERT((std::is_same_v, + nvbench::type_list<>>)); + ASSERT( + (std::is_same_v, + nvbench::type_list< + std::integral_constant>>)); + ASSERT( + (std::is_same_v, + nvbench::type_list< + std::integral_constant, + std::integral_constant, + std::integral_constant>>)); +} + +void test_scoped_enum_type_strings() +{ + using values = nvbench::enum_type_list; + using val_1 = nvbench::tl::get<0, values>; + using val_2 = nvbench::tl::get<1, values>; + using val_3 = nvbench::tl::get<2, values>; + ASSERT((nvbench::type_strings::input_string() == "1")); + ASSERT((nvbench::type_strings::description() == "scoped_enum::val_1")); + ASSERT((nvbench::type_strings::input_string() == "2")); + ASSERT((nvbench::type_strings::description() == "scoped_enum::val_2")); + ASSERT((nvbench::type_strings::input_string() == "3")); + ASSERT((nvbench::type_strings::description() == "scoped_enum::val_3")); +} + +void test_unscoped_enum_type_strings() +{ + using values = nvbench::enum_type_list; + using val_1 = nvbench::tl::get<0, values>; + using val_2 = nvbench::tl::get<1, values>; + using val_3 = nvbench::tl::get<2, values>; + ASSERT((nvbench::type_strings::input_string() == "1")); + ASSERT((nvbench::type_strings::description() == "unscoped_val_1")); + ASSERT((nvbench::type_strings::input_string() == "2")); + ASSERT((nvbench::type_strings::description() == "unscoped_val_2")); + ASSERT((nvbench::type_strings::input_string() == "3")); + ASSERT((nvbench::type_strings::description() == "unscoped_val_3")); +} + +int main() +{ + test_int(); + test_scoped_enum(); + test_unscoped_enum(); + test_scoped_enum_type_strings(); + test_unscoped_enum_type_strings(); +}