diff --git a/cub/cub/device/device_histogram.cuh b/cub/cub/device/device_histogram.cuh index 085df42dc5..679048eef1 100644 --- a/cub/cub/device/device_histogram.cuh +++ b/cub/cub/device/device_histogram.cuh @@ -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) { @@ -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) @@ -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, @@ -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, @@ -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; - 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, @@ -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) @@ -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, @@ -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, @@ -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) { @@ -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) @@ -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, @@ -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, diff --git a/cub/cub/device/dispatch/dispatch_histogram.cuh b/cub/cub/device/dispatch/dispatch_histogram.cuh index 80d4233487..9bf4093513 100644 --- a/cub/cub/device/dispatch/dispatch_histogram.cuh +++ b/cub/cub/device/dispatch/dispatch_histogram.cuh @@ -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; @@ -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, @@ -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, @@ -944,7 +944,7 @@ public: } // Use the search transform op for converting samples to privatized bins - typedef SearchTransform PrivatizedDecodeOpT; + typedef SearchTransform PrivatizedDecodeOpT; // Use the pass-thru transform op for converting privatized bins to output bins typedef PassThruTransform OutputDecodeOpT; @@ -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, @@ -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, @@ -1149,7 +1149,7 @@ public: typedef PassThruTransform PrivatizedDecodeOpT; // Use the search transform op for converting privatized bins to output bins - typedef SearchTransform OutputDecodeOpT; + typedef SearchTransform OutputDecodeOpT; int num_privatized_levels[NUM_ACTIVE_CHANNELS]; PrivatizedDecodeOpT privatized_decode_op[NUM_ACTIVE_CHANNELS]{}; @@ -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, @@ -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, @@ -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, @@ -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, @@ -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,