Use cuda::std::array in histogram APIs (NVIDIA#3973)

Author: bernhardmgruber

cub/benchmarks/bench/histogram/even.cu

@@ -79,11 +79,6 @@ static void even(nvbench::state& state, nvbench::type_list<SampleT, CounterT, Of
   SampleT* d_input      = thrust::raw_pointer_cast(input.data());
   CounterT* d_histogram = thrust::raw_pointer_cast(hist.data());
 
-  CounterT* d_histogram1[1] = {d_histogram};
-  int num_levels1[1]        = {num_levels};
-  SampleT lower_level1[1]   = {lower_level};
-  SampleT upper_level1[1]   = {upper_level};
-
   std::uint8_t* d_temp_storage = nullptr;
   std::size_t temp_storage_bytes{};
 
@@ -100,10 +95,10 @@ static void even(nvbench::state& state, nvbench::type_list<SampleT, CounterT, Of
     d_temp_storage,
     temp_storage_bytes,
     d_input,
-    d_histogram1,
-    num_levels1,
-    lower_level1,
-    upper_level1,
+    {d_histogram},
+    {num_levels},
+    {lower_level},
+    {upper_level},
     num_row_pixels,
     num_rows,
     row_stride_samples,
@@ -118,10 +113,10 @@ static void even(nvbench::state& state, nvbench::type_list<SampleT, CounterT, Of
       d_temp_storage,
       temp_storage_bytes,
       d_input,
-      d_histogram1,
-      num_levels1,
-      lower_level1,
-      upper_level1,
+      {d_histogram},
+      {num_levels},
+      {lower_level},
+      {upper_level},
       num_row_pixels,
       num_rows,
       row_stride_samples,

cub/benchmarks/bench/histogram/multi/even.cu

@@ -89,11 +89,6 @@ static void even(nvbench::state& state, nvbench::type_list<SampleT, CounterT, Of
   CounterT* d_histogram_g = thrust::raw_pointer_cast(hist_g.data());
   CounterT* d_histogram_b = thrust::raw_pointer_cast(hist_b.data());
 
-  CounterT* d_histogram[num_active_channels] = {d_histogram_r, d_histogram_g, d_histogram_b};
-  int num_levels[num_active_channels]        = {num_levels_r, num_levels_g, num_levels_b};
-  SampleT lower_level[num_active_channels]   = {lower_level_r, lower_level_g, lower_level_b};
-  SampleT upper_level[num_active_channels]   = {upper_level_r, upper_level_g, upper_level_b};
-
   std::uint8_t* d_temp_storage = nullptr;
   std::size_t temp_storage_bytes{};
 
@@ -110,10 +105,10 @@ static void even(nvbench::state& state, nvbench::type_list<SampleT, CounterT, Of
     d_temp_storage,
     temp_storage_bytes,
     d_input,
-    d_histogram,
-    num_levels,
-    lower_level,
-    upper_level,
+    {d_histogram_r, d_histogram_g, d_histogram_b},
+    {num_levels_r, num_levels_g, num_levels_b},
+    {lower_level_r, lower_level_g, lower_level_b},
+    {upper_level_r, upper_level_g, upper_level_b},
     num_row_pixels,
     num_rows,
     row_stride_samples,
@@ -128,10 +123,10 @@ static void even(nvbench::state& state, nvbench::type_list<SampleT, CounterT, Of
       d_temp_storage,
       temp_storage_bytes,
       d_input,
-      d_histogram,
-      num_levels,
-      lower_level,
-      upper_level,
+      {d_histogram_r, d_histogram_g, d_histogram_b},
+      {num_levels_r, num_levels_g, num_levels_b},
+      {lower_level_r, lower_level_g, lower_level_b},
+      {upper_level_r, upper_level_g, upper_level_b},
       num_row_pixels,
       num_rows,
       row_stride_samples,

cub/benchmarks/bench/histogram/multi/range.cu

@@ -99,10 +99,6 @@ static void range(nvbench::state& state, nvbench::type_list<SampleT, CounterT, O
   CounterT* d_histogram_g = thrust::raw_pointer_cast(hist_g.data());
   CounterT* d_histogram_b = thrust::raw_pointer_cast(hist_b.data());
 
-  CounterT* d_histogram[num_active_channels] = {d_histogram_r, d_histogram_g, d_histogram_b};
-  int num_levels[num_active_channels]        = {num_levels_r, num_levels_g, num_levels_b};
-  SampleT* d_levels[num_active_channels]     = {d_levels_r, d_levels_g, d_levels_b};
-
   std::uint8_t* d_temp_storage = nullptr;
   std::size_t temp_storage_bytes{};
 
@@ -119,9 +115,9 @@ static void range(nvbench::state& state, nvbench::type_list<SampleT, CounterT, O
     d_temp_storage,
     temp_storage_bytes,
     d_input,
-    d_histogram,
-    num_levels,
-    d_levels,
+    {d_histogram_r, d_histogram_g, d_histogram_b},
+    {num_levels_r, num_levels_g, num_levels_b},
+    {d_levels_r, d_levels_g, d_levels_b},
     num_row_pixels,
     num_rows,
     row_stride_samples,
@@ -136,9 +132,9 @@ static void range(nvbench::state& state, nvbench::type_list<SampleT, CounterT, O
       d_temp_storage,
       temp_storage_bytes,
       d_input,
-      d_histogram,
-      num_levels,
-      d_levels,
+      {d_histogram_r, d_histogram_g, d_histogram_b},
+      {num_levels_r, num_levels_g, num_levels_b},
+      {d_levels_r, d_levels_g, d_levels_b},
       num_row_pixels,
       num_rows,
       row_stride_samples,

cub/benchmarks/bench/histogram/range.cu

@@ -88,10 +88,6 @@ static void range(nvbench::state& state, nvbench::type_list<SampleT, CounterT, O
   SampleT* d_input      = thrust::raw_pointer_cast(input.data());
   CounterT* d_histogram = thrust::raw_pointer_cast(hist.data());
 
-  CounterT* d_histogram1[1] = {d_histogram};
-  int num_levels1[1]        = {num_levels};
-  SampleT* d_levels1[1]     = {d_levels};
-
   std::uint8_t* d_temp_storage = nullptr;
   std::size_t temp_storage_bytes{};
 
@@ -108,9 +104,9 @@ static void range(nvbench::state& state, nvbench::type_list<SampleT, CounterT, O
     d_temp_storage,
     temp_storage_bytes,
     d_input,
-    d_histogram1,
-    num_levels1,
-    d_levels1,
+    {d_histogram},
+    {num_levels},
+    {d_levels},
     num_row_pixels,
     num_rows,
     row_stride_samples,
@@ -125,9 +121,9 @@ static void range(nvbench::state& state, nvbench::type_list<SampleT, CounterT, O
       d_temp_storage,
       temp_storage_bytes,
       d_input,
-      d_histogram1,
-      num_levels1,
-      d_levels1,
+      {d_histogram},
+      {num_levels},
+      {d_levels},
       num_row_pixels,
       num_rows,
       row_stride_samples,