Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Const-qualify histogram pointer input parameters #1762

Merged
merged 1 commit into from
May 23, 2024
Merged

Const-qualify histogram pointer input parameters #1762

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
60 changes: 30 additions & 30 deletions cub/cub/device/device_histogram.cuh
View file
Open in desktop
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
Open in desktop
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
Loading