Skip to content

Commit

Permalink
Const-qualify histogram pointer input parameters
Browse files Browse the repository at this point in the history
  • Loading branch information
@bernhardmgruber
bernhardmgruber committed May 20, 2024
1 parent 0e350f6 commit 5739c0e
Show file tree
Hide file tree
Showing 2 changed files with 55 additions and 55 deletions.
60 changes: 30 additions & 30 deletions cub/cub/device/device_histogram.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -571,9 +571,9 @@ struct DeviceHistogram
size_t& temp_storage_bytes,
SampleIteratorT d_samples,
CounterT* d_histogram[NUM_ACTIVE_CHANNELS],
int num_levels[NUM_ACTIVE_CHANNELS],
LevelT lower_level[NUM_ACTIVE_CHANNELS],
LevelT upper_level[NUM_ACTIVE_CHANNELS],
const int num_levels[NUM_ACTIVE_CHANNELS],
const LevelT lower_level[NUM_ACTIVE_CHANNELS],
const LevelT upper_level[NUM_ACTIVE_CHANNELS],
OffsetT num_pixels,
cudaStream_t stream = 0)
{
Expand Down Expand Up @@ -605,9 +605,9 @@ struct DeviceHistogram
size_t& temp_storage_bytes,
SampleIteratorT d_samples,
CounterT* d_histogram[NUM_ACTIVE_CHANNELS],
int num_levels[NUM_ACTIVE_CHANNELS],
LevelT lower_level[NUM_ACTIVE_CHANNELS],
LevelT upper_level[NUM_ACTIVE_CHANNELS],
const int num_levels[NUM_ACTIVE_CHANNELS],
const LevelT lower_level[NUM_ACTIVE_CHANNELS],
const LevelT upper_level[NUM_ACTIVE_CHANNELS],
OffsetT num_pixels,
cudaStream_t stream,
bool debug_synchronous)
Expand Down Expand Up @@ -789,9 +789,9 @@ struct DeviceHistogram
size_t& temp_storage_bytes,
SampleIteratorT d_samples,
CounterT* d_histogram[NUM_ACTIVE_CHANNELS],
int num_levels[NUM_ACTIVE_CHANNELS],
LevelT lower_level[NUM_ACTIVE_CHANNELS],
LevelT upper_level[NUM_ACTIVE_CHANNELS],
const int num_levels[NUM_ACTIVE_CHANNELS],
const LevelT lower_level[NUM_ACTIVE_CHANNELS],
const LevelT upper_level[NUM_ACTIVE_CHANNELS],
OffsetT num_row_pixels,
OffsetT num_rows,
size_t row_stride_bytes,
Expand Down Expand Up @@ -850,9 +850,9 @@ struct DeviceHistogram
size_t& temp_storage_bytes,
SampleIteratorT d_samples,
CounterT* d_histogram[NUM_ACTIVE_CHANNELS],
int num_levels[NUM_ACTIVE_CHANNELS],
LevelT lower_level[NUM_ACTIVE_CHANNELS],
LevelT upper_level[NUM_ACTIVE_CHANNELS],
const int num_levels[NUM_ACTIVE_CHANNELS],
const LevelT lower_level[NUM_ACTIVE_CHANNELS],
const LevelT upper_level[NUM_ACTIVE_CHANNELS],
OffsetT num_row_pixels,
OffsetT num_rows,
size_t row_stride_bytes,
Expand Down Expand Up @@ -981,16 +981,16 @@ struct DeviceHistogram
SampleIteratorT d_samples,
CounterT* d_histogram,
int num_levels,
LevelT* d_levels,
const LevelT* d_levels,
OffsetT num_samples,
cudaStream_t stream = 0)
{
/// The sample value type of the input iterator
using SampleT = cub::detail::value_t<SampleIteratorT>;

CounterT* d_histogram1[1] = {d_histogram};
int num_levels1[1] = {num_levels};
LevelT* d_levels1[1] = {d_levels};
CounterT* d_histogram1[1] = {d_histogram};
int num_levels1[1] = {num_levels};
const LevelT* d_levels1[1] = {d_levels};

return MultiHistogramRange<1, 1>(
d_temp_storage,
Expand All @@ -1012,7 +1012,7 @@ struct DeviceHistogram
SampleIteratorT d_samples,
CounterT* d_histogram,
int num_levels,
LevelT* d_levels,
const LevelT* d_levels,
OffsetT num_samples,
cudaStream_t stream,
bool debug_synchronous)
Expand Down Expand Up @@ -1142,15 +1142,15 @@ struct DeviceHistogram
SampleIteratorT d_samples,
CounterT* d_histogram,
int num_levels,
LevelT* d_levels,
const LevelT* d_levels,
OffsetT num_row_samples,
OffsetT num_rows,
size_t row_stride_bytes,
cudaStream_t stream = 0)
{
CounterT* d_histogram1[1] = {d_histogram};
int num_levels1[1] = {num_levels};
LevelT* d_levels1[1] = {d_levels};
CounterT* d_histogram1[1] = {d_histogram};
int num_levels1[1] = {num_levels};
const LevelT* d_levels1[1] = {d_levels};

return MultiHistogramRange<1, 1>(
d_temp_storage,
Expand All @@ -1172,7 +1172,7 @@ struct DeviceHistogram
SampleIteratorT d_samples,
CounterT* d_histogram,
int num_levels,
LevelT* d_levels,
const LevelT* d_levels,
OffsetT num_row_samples,
OffsetT num_rows,
size_t row_stride_bytes,
Expand Down Expand Up @@ -1330,8 +1330,8 @@ struct DeviceHistogram
size_t& temp_storage_bytes,
SampleIteratorT d_samples,
CounterT* d_histogram[NUM_ACTIVE_CHANNELS],
int num_levels[NUM_ACTIVE_CHANNELS],
LevelT* d_levels[NUM_ACTIVE_CHANNELS],
const int num_levels[NUM_ACTIVE_CHANNELS],
const LevelT* const d_levels[NUM_ACTIVE_CHANNELS],
OffsetT num_pixels,
cudaStream_t stream = 0)
{
Expand Down Expand Up @@ -1362,8 +1362,8 @@ struct DeviceHistogram
size_t& temp_storage_bytes,
SampleIteratorT d_samples,
CounterT* d_histogram[NUM_ACTIVE_CHANNELS],
int num_levels[NUM_ACTIVE_CHANNELS],
LevelT* d_levels[NUM_ACTIVE_CHANNELS],
const int num_levels[NUM_ACTIVE_CHANNELS],
const LevelT* const d_levels[NUM_ACTIVE_CHANNELS],
OffsetT num_pixels,
cudaStream_t stream,
bool debug_synchronous)
Expand Down Expand Up @@ -1529,8 +1529,8 @@ struct DeviceHistogram
size_t& temp_storage_bytes,
SampleIteratorT d_samples,
CounterT* d_histogram[NUM_ACTIVE_CHANNELS],
int num_levels[NUM_ACTIVE_CHANNELS],
LevelT* d_levels[NUM_ACTIVE_CHANNELS],
const int num_levels[NUM_ACTIVE_CHANNELS],
const LevelT* const d_levels[NUM_ACTIVE_CHANNELS],
OffsetT num_row_pixels,
OffsetT num_rows,
size_t row_stride_bytes,
Expand Down Expand Up @@ -1587,8 +1587,8 @@ struct DeviceHistogram
size_t& temp_storage_bytes,
SampleIteratorT d_samples,
CounterT* d_histogram[NUM_ACTIVE_CHANNELS],
int num_levels[NUM_ACTIVE_CHANNELS],
LevelT* d_levels[NUM_ACTIVE_CHANNELS],
const int num_levels[NUM_ACTIVE_CHANNELS],
const LevelT* const d_levels[NUM_ACTIVE_CHANNELS],
OffsetT num_row_pixels,
OffsetT num_rows,
size_t row_stride_bytes,
Expand Down
50 changes: 25 additions & 25 deletions cub/cub/device/dispatch/dispatch_histogram.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -269,9 +269,9 @@ struct dispatch_histogram
size_t& temp_storage_bytes;
SampleIteratorT d_samples;
CounterT** d_output_histograms;
int* num_privatized_levels;
const int* num_privatized_levels;
PrivatizedDecodeOpT* privatized_decode_op;
int* num_output_levels;
const int* num_output_levels;
OutputDecodeOpT* output_decode_op;
int max_num_output_bins;
OffsetT num_row_pixels;
Expand All @@ -284,9 +284,9 @@ struct dispatch_histogram
size_t& temp_storage_bytes,
SampleIteratorT d_samples,
CounterT* d_output_histograms[NUM_ACTIVE_CHANNELS],
int num_privatized_levels[NUM_ACTIVE_CHANNELS],
const int num_privatized_levels[NUM_ACTIVE_CHANNELS],
PrivatizedDecodeOpT privatized_decode_op[NUM_ACTIVE_CHANNELS],
int num_output_levels[NUM_ACTIVE_CHANNELS],
const int num_output_levels[NUM_ACTIVE_CHANNELS],
OutputDecodeOpT output_decode_op[NUM_ACTIVE_CHANNELS],
int max_num_output_bins,
OffsetT num_row_pixels,
Expand Down Expand Up @@ -922,8 +922,8 @@ public:
size_t& temp_storage_bytes,
SampleIteratorT d_samples,
CounterT* d_output_histograms[NUM_ACTIVE_CHANNELS],
int num_output_levels[NUM_ACTIVE_CHANNELS],
LevelT* d_levels[NUM_ACTIVE_CHANNELS],
const int num_output_levels[NUM_ACTIVE_CHANNELS],
const LevelT* const d_levels[NUM_ACTIVE_CHANNELS],
OffsetT num_row_pixels,
OffsetT num_rows,
OffsetT row_stride_samples,
Expand All @@ -944,7 +944,7 @@ public:
}

// Use the search transform op for converting samples to privatized bins
typedef SearchTransform<LevelT*> PrivatizedDecodeOpT;
typedef SearchTransform<const LevelT*> PrivatizedDecodeOpT;

// Use the pass-thru transform op for converting privatized bins to output bins
typedef PassThruTransform OutputDecodeOpT;
Expand Down Expand Up @@ -1048,7 +1048,7 @@ public:
SampleIteratorT d_samples,
CounterT* d_output_histograms[NUM_ACTIVE_CHANNELS],
int num_output_levels[NUM_ACTIVE_CHANNELS],
LevelT* d_levels[NUM_ACTIVE_CHANNELS],
const LevelT* const d_levels[NUM_ACTIVE_CHANNELS],
OffsetT num_row_pixels,
OffsetT num_rows,
OffsetT row_stride_samples,
Expand Down Expand Up @@ -1124,8 +1124,8 @@ public:
size_t& temp_storage_bytes,
SampleIteratorT d_samples,
CounterT* d_output_histograms[NUM_ACTIVE_CHANNELS],
int num_output_levels[NUM_ACTIVE_CHANNELS],
LevelT* d_levels[NUM_ACTIVE_CHANNELS],
const int num_output_levels[NUM_ACTIVE_CHANNELS],
const LevelT* const d_levels[NUM_ACTIVE_CHANNELS],
OffsetT num_row_pixels,
OffsetT num_rows,
OffsetT row_stride_samples,
Expand All @@ -1149,7 +1149,7 @@ public:
typedef PassThruTransform PrivatizedDecodeOpT;

// Use the search transform op for converting privatized bins to output bins
typedef SearchTransform<LevelT*> OutputDecodeOpT;
typedef SearchTransform<const LevelT*> OutputDecodeOpT;

int num_privatized_levels[NUM_ACTIVE_CHANNELS];
PrivatizedDecodeOpT privatized_decode_op[NUM_ACTIVE_CHANNELS]{};
Expand Down Expand Up @@ -1211,8 +1211,8 @@ public:
size_t& temp_storage_bytes,
SampleIteratorT d_samples,
CounterT* d_output_histograms[NUM_ACTIVE_CHANNELS],
int num_output_levels[NUM_ACTIVE_CHANNELS],
LevelT* d_levels[NUM_ACTIVE_CHANNELS],
const int num_output_levels[NUM_ACTIVE_CHANNELS],
const LevelT* const d_levels[NUM_ACTIVE_CHANNELS],
OffsetT num_row_pixels,
OffsetT num_rows,
OffsetT row_stride_samples,
Expand Down Expand Up @@ -1289,9 +1289,9 @@ public:
size_t& temp_storage_bytes,
SampleIteratorT d_samples,
CounterT* d_output_histograms[NUM_ACTIVE_CHANNELS],
int num_output_levels[NUM_ACTIVE_CHANNELS],
LevelT lower_level[NUM_ACTIVE_CHANNELS],
LevelT upper_level[NUM_ACTIVE_CHANNELS],
const int num_output_levels[NUM_ACTIVE_CHANNELS],
const LevelT lower_level[NUM_ACTIVE_CHANNELS],
const LevelT upper_level[NUM_ACTIVE_CHANNELS],
OffsetT num_row_pixels,
OffsetT num_rows,
OffsetT row_stride_samples,
Expand Down Expand Up @@ -1427,9 +1427,9 @@ public:
size_t& temp_storage_bytes,
SampleIteratorT d_samples,
CounterT* d_output_histograms[NUM_ACTIVE_CHANNELS],
int num_output_levels[NUM_ACTIVE_CHANNELS],
LevelT lower_level[NUM_ACTIVE_CHANNELS],
LevelT upper_level[NUM_ACTIVE_CHANNELS],
const int num_output_levels[NUM_ACTIVE_CHANNELS],
const LevelT lower_level[NUM_ACTIVE_CHANNELS],
const LevelT upper_level[NUM_ACTIVE_CHANNELS],
OffsetT num_row_pixels,
OffsetT num_rows,
OffsetT row_stride_samples,
Expand Down Expand Up @@ -1508,9 +1508,9 @@ public:
size_t& temp_storage_bytes,
SampleIteratorT d_samples,
CounterT* d_output_histograms[NUM_ACTIVE_CHANNELS],
int num_output_levels[NUM_ACTIVE_CHANNELS],
LevelT lower_level[NUM_ACTIVE_CHANNELS],
LevelT upper_level[NUM_ACTIVE_CHANNELS],
const int num_output_levels[NUM_ACTIVE_CHANNELS],
const LevelT lower_level[NUM_ACTIVE_CHANNELS],
const LevelT upper_level[NUM_ACTIVE_CHANNELS],
OffsetT num_row_pixels,
OffsetT num_rows,
OffsetT row_stride_samples,
Expand Down Expand Up @@ -1596,9 +1596,9 @@ public:
size_t& temp_storage_bytes,
SampleIteratorT d_samples,
CounterT* d_output_histograms[NUM_ACTIVE_CHANNELS],
int num_output_levels[NUM_ACTIVE_CHANNELS],
LevelT lower_level[NUM_ACTIVE_CHANNELS],
LevelT upper_level[NUM_ACTIVE_CHANNELS],
const int num_output_levels[NUM_ACTIVE_CHANNELS],
const LevelT lower_level[NUM_ACTIVE_CHANNELS],
const LevelT upper_level[NUM_ACTIVE_CHANNELS],
OffsetT num_row_pixels,
OffsetT num_rows,
OffsetT row_stride_samples,
Expand Down

0 comments on commit 5739c0e

Please sign in to comment.