Cleanup bench macro annotation

gevtushenko · May 10, 2023 · b8837dc · b8837dc
1 parent f200d0c
commit b8837dc
Show file tree

Hide file tree

Showing 11 changed files with 76 additions and 76 deletions.
diff --git a/benchmarks/bench/merge_sort/keys.cu b/benchmarks/bench/merge_sort/keys.cu
@@ -9,7 +9,7 @@
 
 #ifndef TUNE_BASE
 #define TUNE_THREADS_PER_BLOCK (1 << TUNE_THREADS_PER_BLOCK_POW2)
-#endif
+#endif // TUNE_BASE
 
 using value_t = cub::NullType;
 
@@ -18,18 +18,18 @@ using value_t = cub::NullType;
 #if TUNE_TRANSPOSE == 0
 #define TUNE_LOAD_ALGORITHM cub::BLOCK_LOAD_DIRECT
 #define TUNE_STORE_ALGORITHM cub::BLOCK_STORE_DIRECT
-#else
+#else // TUNE_TRANSPOSE == 1
 #define TUNE_LOAD_ALGORITHM cub::BLOCK_LOAD_WARP_TRANSPOSE
 #define TUNE_STORE_ALGORITHM cub::BLOCK_STORE_WARP_TRANSPOSE
-#endif
+#endif // TUNE_TRANSPOSE
 
 #if TUNE_LOAD == 0
 #define TUNE_LOAD_MODIFIER cub::LOAD_DEFAULT
 #elif TUNE_LOAD == 1
 #define TUNE_LOAD_MODIFIER cub::LOAD_LDG
-#else
+#else // TUNE_LOAD == 2
 #define TUNE_LOAD_MODIFIER cub::LOAD_CA
-#endif
+#endif // TUNE_LOAD 
 
 template <typename KeyT>
 struct policy_hub_t
@@ -46,7 +46,7 @@ struct policy_hub_t
 
   using MaxPolicy = policy_t;
 };
-#endif
+#endif // !TUNE_BASE
 
 template <typename T, typename OffsetT>
 void merge_sort_keys(nvbench::state &state, nvbench::type_list<T, OffsetT>)
@@ -69,10 +69,10 @@ void merge_sort_keys(nvbench::state &state, nvbench::type_list<T, OffsetT>)
                                             offset_t,
                                             compare_op_t,
                                             policy_t>;
-#else
+#else // TUNE_BASE
   using dispatch_t = cub::
     DispatchMergeSort<key_input_it_t, value_input_it_t, key_it_t, value_it_t, offset_t, compare_op_t>;
-#endif
+#endif // TUNE_BASE
 
   // Retrieve axis parameters
   const auto elements       = static_cast<std::size_t>(state.get_int64("Elements{io}"));

diff --git a/benchmarks/bench/merge_sort/pairs.cu b/benchmarks/bench/merge_sort/pairs.cu
@@ -15,18 +15,18 @@
 #if TUNE_TRANSPOSE == 0
 #define TUNE_LOAD_ALGORITHM cub::BLOCK_LOAD_DIRECT
 #define TUNE_STORE_ALGORITHM cub::BLOCK_STORE_DIRECT
-#else
+#else // TUNE_TRANSPOSE == 1
 #define TUNE_LOAD_ALGORITHM cub::BLOCK_LOAD_WARP_TRANSPOSE
 #define TUNE_STORE_ALGORITHM cub::BLOCK_STORE_WARP_TRANSPOSE
-#endif
+#endif // TUNE_TRANSPOSE 
 
 #if TUNE_LOAD == 0
 #define TUNE_LOAD_MODIFIER cub::LOAD_DEFAULT
 #elif TUNE_LOAD == 1
 #define TUNE_LOAD_MODIFIER cub::LOAD_LDG
-#else
+#else // TUNE_LOAD == 2
 #define TUNE_LOAD_MODIFIER cub::LOAD_CA
-#endif
+#endif // TUNE_LOAD
 
 template <typename KeyT>
 struct policy_hub_t
@@ -43,7 +43,7 @@ struct policy_hub_t
 
   using MaxPolicy = policy_t;
 };
-#endif
+#endif // TUNE_BASE
 
 template <typename KeyT, typename ValueT, typename OffsetT>
 void merge_sort_keys(nvbench::state &state, nvbench::type_list<KeyT, ValueT, OffsetT>)
@@ -66,10 +66,10 @@ void merge_sort_keys(nvbench::state &state, nvbench::type_list<KeyT, ValueT, Off
                                             offset_t,
                                             compare_op_t,
                                             policy_t>;
-#else
+#else // TUNE_BASE
   using dispatch_t = cub::
     DispatchMergeSort<key_input_it_t, value_input_it_t, key_it_t, value_it_t, offset_t, compare_op_t>;
-#endif
+#endif // TUNE_BASE
 
   // Retrieve axis parameters
   const auto elements       = static_cast<std::size_t>(state.get_int64("Elements{io}"));
@@ -124,15 +124,15 @@ void merge_sort_keys(nvbench::state &state, nvbench::type_list<KeyT, ValueT, Off
 
 #ifdef TUNE_KeyT
 using key_types = nvbench::type_list<TUNE_KeyT>;
-#else
+#else // !defined(TUNE_KeyT)
 using key_types = all_types;
-#endif
+#endif // TUNE_KeyT
 
 #ifdef TUNE_ValueT
 using value_types = nvbench::type_list<TUNE_ValueT>;
-#else
+#else // !defined(TUNE_ValueT)
 using value_types = nvbench::type_list<int8_t, int16_t, int32_t, int64_t, int128_t>;
-#endif
+#endif // TUNE_ValueT
 
 NVBENCH_BENCH_TYPES(merge_sort_keys, NVBENCH_TYPE_AXES(key_types, value_types, offset_types))
   .set_name("cub::DeviceMergeSort::SortPairs")

diff --git a/benchmarks/bench/partition/flagged.cu b/benchmarks/bench/partition/flagged.cu
@@ -14,15 +14,15 @@ constexpr bool may_alias = false;
 #if !TUNE_BASE
 #if TUNE_TRANSPOSE == 0
 #define TUNE_LOAD_ALGORITHM cub::BLOCK_LOAD_DIRECT
-#else
+#else // TUNE_TRANSPOSE == 1
 #define TUNE_LOAD_ALGORITHM cub::BLOCK_LOAD_WARP_TRANSPOSE
-#endif
+#endif // TUNE_TRANSPOSE 
 
 #if TUNE_LOAD == 0
 #define TUNE_LOAD_MODIFIER cub::LOAD_DEFAULT
-#else
+#else // TUNE_LOAD == 1
 #define TUNE_LOAD_MODIFIER cub::LOAD_CA
-#endif
+#endif // TUNE_LOAD 
 
 template <typename InputT>
 struct policy_hub_t
@@ -44,7 +44,7 @@ struct policy_hub_t
 
   using MaxPolicy = policy_t;
 };
-#endif
+#endif // TUNE_BASE
 
 template <typename T, typename OffsetT>
 void partition(nvbench::state &state, nvbench::type_list<T, OffsetT>)
@@ -57,7 +57,7 @@ void partition(nvbench::state &state, nvbench::type_list<T, OffsetT>)
   using equality_op_t = cub::NullType;
   using offset_t = OffsetT;
 
-#if !TUNE_BASE
+  #if !TUNE_BASE
   using policy_t = policy_hub_t<T>;
   using dispatch_t = cub::DispatchSelectIf<input_it_t,
                                            flag_it_t,
@@ -69,7 +69,7 @@ void partition(nvbench::state &state, nvbench::type_list<T, OffsetT>)
                                            keep_rejects,
                                            may_alias,
                                            policy_t>;
-#else
+  #else // TUNE_BASE
   using dispatch_t = cub::DispatchSelectIf<input_it_t,
                                            flag_it_t,
                                            output_it_t,
@@ -79,7 +79,7 @@ void partition(nvbench::state &state, nvbench::type_list<T, OffsetT>)
                                            offset_t,
                                            keep_rejects,
                                            may_alias>;
-#endif
+  #endif // TUNE_BASE
 
   // Retrieve axis parameters
   const auto elements = static_cast<std::size_t>(state.get_int64("Elements{io}"));

diff --git a/benchmarks/bench/partition/if.cu b/benchmarks/bench/partition/if.cu
@@ -14,15 +14,15 @@ constexpr bool may_alias = false;
 #if !TUNE_BASE
 #if TUNE_TRANSPOSE == 0
 #define TUNE_LOAD_ALGORITHM cub::BLOCK_LOAD_DIRECT
-#else
+#else // TUNE_TRANSPOSE == 1
 #define TUNE_LOAD_ALGORITHM cub::BLOCK_LOAD_WARP_TRANSPOSE
-#endif
+#endif // TUNE_TRANSPOSE
 
 #if TUNE_LOAD == 0
 #define TUNE_LOAD_MODIFIER cub::LOAD_DEFAULT
-#else
+#else // TUNE_LOAD == 1
 #define TUNE_LOAD_MODIFIER cub::LOAD_CA
-#endif
+#endif // TUNE_LOAD
 
 template <typename InputT>
 struct policy_hub_t
@@ -44,7 +44,7 @@ struct policy_hub_t
 
   using MaxPolicy = policy_t;
 };
-#endif
+#endif // !TUNE_BASE
 
 template <class T>
 struct less_then_t 
@@ -79,7 +79,7 @@ void partition(nvbench::state &state, nvbench::type_list<T, OffsetT>)
   using equality_op_t = cub::NullType;
   using offset_t = OffsetT;
 
-#if !TUNE_BASE
+  #if !TUNE_BASE
   using policy_t = policy_hub_t<T>;
   using dispatch_t = cub::DispatchSelectIf<input_it_t,
                                            flag_it_t,
@@ -91,7 +91,7 @@ void partition(nvbench::state &state, nvbench::type_list<T, OffsetT>)
                                            keep_rejects,
                                            may_alias,
                                            policy_t>;
-#else
+  #else // TUNE_BASE
   using dispatch_t = cub::DispatchSelectIf<input_it_t,
                                            flag_it_t,
                                            output_it_t,
@@ -101,7 +101,7 @@ void partition(nvbench::state &state, nvbench::type_list<T, OffsetT>)
                                            offset_t,
                                            keep_rejects,
                                            may_alias>;
-#endif
+  #endif // !TUNE_BASE
 
   // Retrieve axis parameters
   const auto elements = static_cast<std::size_t>(state.get_int64("Elements{io}"));

diff --git a/benchmarks/bench/radix_sort/keys.cu b/benchmarks/bench/radix_sort/keys.cu
@@ -95,13 +95,13 @@ constexpr bool fits_in_default_shared_memory()
 {
   return max_temp_storage_size<KeyT, ValueT, OffsetT>() < 48 * 1024;
 }
-#else
+#else // TUNE_BASE
 template <typename, typename, typename>
 constexpr bool fits_in_default_shared_memory()
 {
   return true;
 }
-#endif
+#endif // TUNE_BASE
 
 template <typename T, typename OffsetT>
 void radix_sort_keys(std::integral_constant<bool, true>,
@@ -114,9 +114,9 @@ void radix_sort_keys(std::integral_constant<bool, true>,
 #if !TUNE_BASE
   using policy_t   = policy_hub_t<key_t, value_t, offset_t>;
   using dispatch_t = cub::DispatchRadixSort<is_descending, key_t, value_t, offset_t, policy_t>;
-#else
+#else // TUNE_BASE
   using dispatch_t = cub::DispatchRadixSort<is_descending, key_t, value_t, offset_t>;
-#endif
+#endif // TUNE_BASE
 
   const int begin_bit = 0;
   const int end_bit   = sizeof(key_t) * 8;

diff --git a/benchmarks/bench/radix_sort/pairs.cu b/benchmarks/bench/radix_sort/pairs.cu
@@ -93,13 +93,13 @@ constexpr bool fits_in_default_shared_memory()
 {
   return max_temp_storage_size<KeyT, ValueT, OffsetT>() < 48 * 1024;
 }
-#else
+#else // TUNE_BASE
 template <typename, typename, typename>
 constexpr bool fits_in_default_shared_memory()
 {
   return true;
 }
-#endif
+#endif // TUNE_BASE
 
 template <typename KeyT, typename ValueT, typename OffsetT>
 void radix_sort_values(std::integral_constant<bool, true>,
@@ -113,9 +113,9 @@ void radix_sort_values(std::integral_constant<bool, true>,
 #if !TUNE_BASE
   using policy_t   = policy_hub_t<key_t, value_t, offset_t>;
   using dispatch_t = cub::DispatchRadixSort<is_descending, key_t, value_t, offset_t, policy_t>;
-#else
+#else // TUNE_BASE
   using dispatch_t = cub::DispatchRadixSort<is_descending, key_t, value_t, offset_t>;
-#endif
+#endif // TUNE_BASE
 
   const int begin_bit = 0;
   const int end_bit   = sizeof(key_t) * 8;
@@ -201,15 +201,15 @@ void radix_sort_values(nvbench::state &state, nvbench::type_list<KeyT, ValueT, O
 
 #ifdef TUNE_KeyT
 using key_types = nvbench::type_list<TUNE_KeyT>;
-#else
+#else // !defined(TUNE_KeyT) 
 using key_types = fundamental_types;
-#endif
+#endif // TUNE_KeyT
 
 #ifdef TUNE_ValueT
 using value_types = nvbench::type_list<TUNE_ValueT>;
-#else
+#else // !defined(Tune_ValueT)
 using value_types = nvbench::type_list<int8_t, int16_t, int32_t, int64_t, int128_t>;
-#endif
+#endif // TUNE_ValueT
 
 NVBENCH_BENCH_TYPES(radix_sort_values, NVBENCH_TYPE_AXES(key_types, value_types, offset_types))
   .set_name("cub::DeviceRadixSort::SortPairs")

diff --git a/benchmarks/bench/reduce/base.cuh b/benchmarks/bench/reduce/base.cuh
@@ -30,7 +30,7 @@ struct policy_hub_t
 
   using MaxPolicy = policy_t;
 };
-#endif
+#endif // !TUNE_BASE
 
 template <typename T, typename OffsetT>
 void reduce(nvbench::state &state, nvbench::type_list<T, OffsetT>)
@@ -45,9 +45,9 @@ void reduce(nvbench::state &state, nvbench::type_list<T, OffsetT>)
   using policy_t = policy_hub_t<accum_t, offset_t>;
   using dispatch_t =
     cub::DispatchReduce<input_it_t, output_it_t, offset_t, op_t, init_t, accum_t, policy_t>;
-#else
+#else // TUNE_BASE
   using dispatch_t = cub::DispatchReduce<input_it_t, output_it_t, offset_t, op_t, init_t, accum_t>;
-#endif
+#endif // TUNE_BASE
 
   // Retrieve axis parameters
   const auto elements = static_cast<std::size_t>(state.get_int64("Elements{io}"));

diff --git a/benchmarks/bench/scan/exclusive/base.cuh b/benchmarks/bench/scan/exclusive/base.cuh
@@ -47,13 +47,13 @@ constexpr bool fits_in_default_shared_memory()
 {
   return max_temp_storage_size<T, OffsetT>() < 48 * 1024;
 }
-#else
+#else // TUNE_BASE
 template <typename T, typename OffsetT>
 constexpr bool fits_in_default_shared_memory()
 {
   return true;
 }
-#endif
+#endif // TUNE_BASE
 
 template <typename T, typename OffsetT>
 static void basic(std::integral_constant<bool, true>,

diff --git a/benchmarks/bench/scan/exclusive/by_key.cu b/benchmarks/bench/scan/exclusive/by_key.cu
@@ -25,7 +25,7 @@ struct policy_hub_t
 
   using MaxPolicy = policy_t;
 };
-#endif
+#endif // !TUNE_BASE
 
 template <typename KeyT, typename ValueT, typename OffsetT>
 static void scan(nvbench::state &state, nvbench::type_list<KeyT, ValueT, OffsetT>)
@@ -39,7 +39,7 @@ static void scan(nvbench::state &state, nvbench::type_list<KeyT, ValueT, OffsetT
   using equality_op_t   = cub::Equality;
   using offset_t        = OffsetT;
 
-#if !TUNE_BASE
+  #if !TUNE_BASE
   using policy_t   = policy_hub_t;
   using dispatch_t = cub::DispatchScanByKey<key_input_it_t,
                                             val_input_it_t,
@@ -50,7 +50,7 @@ static void scan(nvbench::state &state, nvbench::type_list<KeyT, ValueT, OffsetT
                                             offset_t,
                                             accum_t,
                                             policy_t>;
-#else
+  #else // TUNE_BASE
   using dispatch_t = cub::DispatchScanByKey<key_input_it_t,
                                             val_input_it_t,
                                             val_output_it_t,
@@ -59,7 +59,7 @@ static void scan(nvbench::state &state, nvbench::type_list<KeyT, ValueT, OffsetT
                                             init_value_t,
                                             offset_t,
                                             accum_t>;
-#endif
+  #endif // TUNE_BASE
 
   const auto elements = static_cast<std::size_t>(state.get_int64("Elements{io}"));
 
@@ -112,15 +112,15 @@ using some_offset_types = nvbench::type_list<nvbench::int32_t>;
 
 #ifdef TUNE_KeyT
 using key_types = nvbench::type_list<TUNE_KeyT>;
-#else
+#else // !defined(TUNE_KeyT)
 using key_types = all_types;
-#endif
+#endif // TUNE_KeyT
 
 #ifdef TUNE_ValueT
 using value_types = nvbench::type_list<TUNE_ValueT>;
-#else
+#else // !defined(TUNE_ValueT)
 using value_types = nvbench::type_list<int8_t, int16_t, int32_t, int64_t, int128_t>;
-#endif
+#endif // TUNE_ValueT
 
 NVBENCH_BENCH_TYPES(scan, NVBENCH_TYPE_AXES(key_types, value_types, some_offset_types))
   .set_name("cub::DeviceScan::ExclusiveSumByKey")