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

@@ -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)

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

examples/CMakeLists.txt

@@ -1,5 +1,6 @@
 set(example_srcs
   axes.cu
+  enums.cu
   exec_tag_sync.cu
   exec_tag_timer.cu
   skip.cu

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 <nvbench/nvbench.cuh>
+
+#include <nvbench/test_kernels.cuh>
+
+// 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<MyEnum>(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<nvbench::int64_t>(MyEnum::ValueA),
+                   static_cast<nvbench::int64_t>(MyEnum::ValueB),
+                   static_cast<nvbench::int64_t>(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<Enum, Value>` 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 <MyEnum EnumValue>
+void compile_time_enum_sweep(
+  nvbench::state &state,
+  nvbench::type_list<std::integral_constant<MyEnum, EnumValue>>)
+{
+  // 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<MyEnum, MyEnum::ValueA, MyEnum::ValueB, MyEnum::ValueC>;
+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<int,0>)
+//  * `16` (struct std::integral_constant<int,16>)
+//  * `4096` (struct std::integral_constant<int,4096>)
+//  * `-12` (struct std::integral_constant<int,-12>)
+// ```
+template <nvbench::int32_t IntValue>
+void compile_time_int_sweep(
+  nvbench::state &state,
+  nvbench::type_list<std::integral_constant<nvbench::int32_t, IntValue>>)
+{
+  // 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::int32_t, 0, 16, 4096, -12>;
+NVBENCH_BENCH_TYPES(compile_time_int_sweep, NVBENCH_TYPE_AXES(MyInts))
+  .set_type_axes_names({"SomeInts"});

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 <nvbench/type_list.cuh>
+#include <nvbench/type_strings.cuh>
+
+#include <type_traits>
+
+namespace nvbench
+{
+
+/*!
+ * \brief Helper utility that generates a `type_list` of
+ * `std::integral_constant`s.
+ *
+ * \relatesalso NVBENCH_DECLARE_ENUM_TYPE_STRINGS
+ */
+template <typename T, T... Ts>
+using enum_type_list = nvbench::type_list<std::integral_constant<T, Ts>...>;
+} // 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<std::integral_constant<T, Value>>`.
+ *
+ * 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 <T Value>                                                           \
+  struct type_strings<std::integral_constant<T, Value>>                        \
+  {                                                                            \
+    static std::string input_string() { return input_generator(Value); }       \
+    static std::string description() { return description_generator(Value); }  \
+  };                                                                           \
+  }

nvbench/nvbench.cuh

@@ -27,6 +27,7 @@
 #include <nvbench/cuda_call.cuh>
 #include <nvbench/cuda_stream.cuh>
 #include <nvbench/cuda_timer.cuh>
+#include <nvbench/enum_type_list.cuh>
 #include <nvbench/exec_tag.cuh>
 #include <nvbench/launch.cuh>
 #include <nvbench/main.cuh>

nvbench/type_strings.cu

@@ -18,6 +18,8 @@
 
 #include <nvbench/type_strings.cuh>
 
+#include <fmt/format.h>
+
 #include <string>
 
 #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

nvbench/type_strings.cuh

@@ -21,6 +21,7 @@
 #include <nvbench/types.cuh>
 
 #include <string>
+#include <type_traits>
 #include <typeinfo>
 
 namespace nvbench
@@ -49,8 +50,26 @@ struct type_strings
   static std::string description() { return {}; }
 };
 
+template <typename T, T Value>
+struct type_strings<std::integral_constant<T, Value>>
+{
+  // 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<std::integral_constant<T, Value>>();
+  }
+};
+
 } // namespace nvbench
 
+/*!
+ * Declare an `input_string` and `description` to use with a specific `type`.
+ */
 #define NVBENCH_DECLARE_TYPE_STRINGS(Type, InputString, Description)           \
   namespace nvbench                                                            \
   {                                                                            \

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

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 <nvbench/enum_type_list.cuh>
+
+#include "test_asserts.cuh"
+
+#include <fmt/format.h>
+
+#include <type_traits>
+
+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::enum_type_list<int>, nvbench::type_list<>>));
+  ASSERT((std::is_same_v<nvbench::enum_type_list<int, 0>,
+                         nvbench::type_list<std::integral_constant<int, 0>>>));
+  ASSERT((std::is_same_v<nvbench::enum_type_list<int, 0, 1, 2, 3, 4>,
+                         nvbench::type_list<std::integral_constant<int, 0>,
+                                            std::integral_constant<int, 1>,
+                                            std::integral_constant<int, 2>,
+                                            std::integral_constant<int, 3>,
+                                            std::integral_constant<int, 4>>>));
+}
+
+void test_scoped_enum()
+{
+  ASSERT((
+    std::is_same_v<nvbench::enum_type_list<scoped_enum>, nvbench::type_list<>>));
+  ASSERT((
+    std::is_same_v<nvbench::enum_type_list<scoped_enum, scoped_enum::val_1>,
+                   nvbench::type_list<
+                     std::integral_constant<scoped_enum, scoped_enum::val_1>>>));
+  ASSERT((
+    std::is_same_v<nvbench::enum_type_list<scoped_enum,
+                                           scoped_enum::val_1,
+                                           scoped_enum::val_2,
+                                           scoped_enum::val_3>,
+                   nvbench::type_list<
+                     std::integral_constant<scoped_enum, scoped_enum::val_1>,
+                     std::integral_constant<scoped_enum, scoped_enum::val_2>,
+                     std::integral_constant<scoped_enum, scoped_enum::val_3>>>));
+}
+
+void test_unscoped_enum()
+{
+  ASSERT((std::is_same_v<nvbench::enum_type_list<unscoped_enum>,
+                         nvbench::type_list<>>));
+  ASSERT(
+    (std::is_same_v<nvbench::enum_type_list<unscoped_enum, unscoped_val_1>,
+                    nvbench::type_list<
+                      std::integral_constant<unscoped_enum, unscoped_val_1>>>));
+  ASSERT(
+    (std::is_same_v<nvbench::enum_type_list<unscoped_enum,
+                                            unscoped_val_1,
+                                            unscoped_val_2,
+                                            unscoped_val_3>,
+                    nvbench::type_list<
+                      std::integral_constant<unscoped_enum, unscoped_val_1>,
+                      std::integral_constant<unscoped_enum, unscoped_val_2>,
+                      std::integral_constant<unscoped_enum, unscoped_val_3>>>));
+}
+
+void test_scoped_enum_type_strings()
+{
+  using values = nvbench::enum_type_list<scoped_enum,
+                                         scoped_enum::val_1,
+                                         scoped_enum::val_2,
+                                         scoped_enum::val_3>;
+  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<val_1>::input_string() == "1"));
+  ASSERT((nvbench::type_strings<val_1>::description() == "scoped_enum::val_1"));
+  ASSERT((nvbench::type_strings<val_2>::input_string() == "2"));
+  ASSERT((nvbench::type_strings<val_2>::description() == "scoped_enum::val_2"));
+  ASSERT((nvbench::type_strings<val_3>::input_string() == "3"));
+  ASSERT((nvbench::type_strings<val_3>::description() == "scoped_enum::val_3"));
+}
+
+void test_unscoped_enum_type_strings()
+{
+  using values = nvbench::enum_type_list<unscoped_enum,
+                                         unscoped_enum::unscoped_val_1,
+                                         unscoped_enum::unscoped_val_2,
+                                         unscoped_enum::unscoped_val_3>;
+  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<val_1>::input_string() == "1"));
+  ASSERT((nvbench::type_strings<val_1>::description() == "unscoped_val_1"));
+  ASSERT((nvbench::type_strings<val_2>::input_string() == "2"));
+  ASSERT((nvbench::type_strings<val_2>::description() == "unscoped_val_2"));
+  ASSERT((nvbench::type_strings<val_3>::input_string() == "3"));
+  ASSERT((nvbench::type_strings<val_3>::description() == "unscoped_val_3"));
+}
+
+int main()
+{
+  test_int();
+  test_scoped_enum();
+  test_unscoped_enum();
+  test_scoped_enum_type_strings();
+  test_unscoped_enum_type_strings();
+}