@@ -176,16 +176,20 @@ NVBENCH_BENCH_TYPES(copy_type_conversion_sweep, NVBENCH_TYPE_AXES(ctcs_types, ct
176176// Passing list of typenames and `enum_type_list` to build cartesian product
177177// of typenames and integral constants
178178
179+ // define constant wrapper helper type
180+ template <auto V, typename T = decltype (V)>
181+ using cw_t = std::integral_constant<T, V>;
182+
179183template <typename ValueT, unsigned BLOCK_DIM>
180184void copy_type_and_block_size_sweep (nvbench::state &state,
181- nvbench::type_list<ValueT, nvbench::enum_type <BLOCK_DIM>>)
185+ nvbench::type_list<ValueT, cw_t <BLOCK_DIM>>)
182186{
183187 const std::size_t nelems = 256 * 1024 * 1024 / sizeof (ValueT);
184188 ValueT fill_value{42 };
185189 thrust::device_vector<ValueT> inp (nelems, fill_value);
186190 thrust::device_vector<ValueT> out (nelems, ValueT{});
187191
188- const auto gridSize = cuda::ceil_div (nelems, BLOCK_DIM) ;
192+ const auto gridSize = (nelems + BLOCK_DIM - 1 ) / BLOCK_DIM ;
189193
190194 const ValueT *inp_p = thrust::raw_pointer_cast (inp.data ());
191195 ValueT *out_p = thrust::raw_pointer_cast (out.data ());
@@ -199,6 +203,9 @@ void copy_type_and_block_size_sweep(nvbench::state &state,
199203 });
200204}
201205
202- using block_sizes = nvbench::enum_type_list<64u , 128u , 196u , 256u , 320u , 512u >;
206+ template <auto ... V>
207+ using cw_list = nvbench::type_list<cw_t <V>...>;
208+
209+ using block_sizes = cw_list<64u , 128u , 196u , 256u , 320u , 512u >;
203210NVBENCH_BENCH_TYPES (copy_type_and_block_size_sweep, NVBENCH_TYPE_AXES(ctcs_types, block_sizes))
204211 .set_type_axes_names({" Type" , " BlockSize" });
0 commit comments