Histogram#
Configuring the kernel#
-
template<class HistogramConfig, unsigned int MaxGridSize = 1024, unsigned int SharedImplMaxBins = 2048, unsigned int SharedImplHistograms = 3>
struct histogram_config : public rocprim::detail::histogram_config_params# Configuration of device-level histogram operation.
- Template Parameters:
HistogramConfig – - configuration of histogram kernel. Must be
kernel_config
.MaxGridSize – - maximum number of blocks to launch.
SharedImplMaxBins – - maximum total number of bins for all active channels for the shared memory histogram implementation (samples -> shared memory bins -> global memory bins), when exceeded the global memory implementation is used (samples -> global memory bins).
SharedImplHistograms – - number of histograms in the shared memory to reduce bank conflicts for atomic operations with narrow sample distributions. Sweetspot for 9xx and 10xx is 3.
Subclassed by rocprim::detail::default_histogram_config< arch, value_type, channels, active_channels, enable >
histogram_even#
-
template<class Config = default_config, class SampleIterator, class Counter, class Level>
inline hipError_t rocprim::histogram_even(void *temporary_storage, size_t &storage_size, SampleIterator samples, unsigned int size, Counter *histogram, unsigned int levels, Level lower_level, Level upper_level, hipStream_t stream = 0, bool debug_synchronous = false)# Computes a histogram from a sequence of samples using equal-width bins.
The number of histogram bins is (
levels
- 1).Bins are evenly-segmented and include the same width of sample values: (
upper_level
-lower_level
) / (levels
- 1).Returns the required size of
temporary_storage
instorage_size
iftemporary_storage
in a null pointer.
- Example
In this example a device-level histogram of 5 bins is computed on an array of float samples.
#include <rocprim/rocprim.hpp> // Prepare input and output (declare pointers, allocate device memory etc.) unsigned int size; // e.g., 8 float * samples; // e.g., [-10.0, 0.3, 9.5, 8.1, 1.5, 1.9, 100.0, 5.1] int * histogram; // empty array of at least 5 elements unsigned int levels; // e.g., 6 (for 5 bins) float lower_level; // e.g., 0.0 float upper_level; // e.g., 10.0 size_t temporary_storage_size_bytes; void * temporary_storage_ptr = nullptr; // Get required size of the temporary storage rocprim::histogram_even( temporary_storage_ptr, temporary_storage_size_bytes, samples, size, histogram, levels, lower_level, upper_level ); // allocate temporary storage hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes); // compute histogram rocprim::histogram_even( temporary_storage_ptr, temporary_storage_size_bytes, samples, size, histogram, levels, lower_level, upper_level ); // histogram: [3, 0, 1, 0, 2]
- Template Parameters:
Config – - [optional] configuration of the primitive. It has to be
histogram_config
or a class derived from it.SampleIterator – - random-access iterator type of the input range. Must meet the requirements of a C++ InputIterator concept. It can be a simple pointer type.
Counter – - integer type for histogram bin counters.
Level – - type of histogram boundaries (levels)
- Parameters:
temporary_storage – [in] - pointer to a device-accessible temporary storage. When a null pointer is passed, the required allocation size (in bytes) is written to
storage_size
and function returns without performing the reduction operation.storage_size – [inout] - reference to a size (in bytes) of
temporary_storage
.samples – [in] - iterator to the first element in the range of input samples.
size – [in] - number of elements in the samples range.
histogram – [out] - pointer to the first element in the histogram range.
levels – [in] - number of boundaries (levels) for histogram bins.
lower_level – [in] - lower sample value bound (inclusive) for the first histogram bin.
upper_level – [in] - upper sample value bound (exclusive) for the last histogram bin.
stream – [in] - [optional] HIP stream object. Default is
0
(default stream).debug_synchronous – [in] - [optional] If true, synchronization after every kernel launch is forced in order to check for errors. Default value is
false
.
- Returns:
hipSuccess
(0
) after successful histogram operation; otherwise a HIP runtime error of typehipError_t
.
-
template<class Config = default_config, class SampleIterator, class Counter, class Level>
inline hipError_t rocprim::histogram_even(void *temporary_storage, size_t &storage_size, SampleIterator samples, unsigned int columns, unsigned int rows, size_t row_stride_bytes, Counter *histogram, unsigned int levels, Level lower_level, Level upper_level, hipStream_t stream = 0, bool debug_synchronous = false)# Computes a histogram from a two-dimensional region of samples using equal-width bins.
The two-dimensional region of interest within
samples
can be specified using thecolumns
,rows
androw_stride_bytes
parameters.The row stride must be a whole multiple of the sample data type size, i.e.,
(row_stride_bytes % sizeof(std::iterator_traits<SampleIterator>::value_type)) == 0
.The number of histogram bins is (
levels
- 1).Bins are evenly-segmented and include the same width of sample values: (
upper_level
-lower_level
) / (levels
- 1).Returns the required size of
temporary_storage
instorage_size
iftemporary_storage
in a null pointer.
- Example
In this example a device-level histogram of 5 bins is computed on an array of float samples.
#include <rocprim/rocprim.hpp> // Prepare input and output (declare pointers, allocate device memory etc.) unsigned int columns; // e.g., 4 unsigned int rows; // e.g., 2 size_t row_stride_bytes; // e.g., 6 * sizeof(float) float * samples; // e.g., [-10.0, 0.3, 9.5, 8.1, -, -, 1.5, 1.9, 100.0, 5.1, -, -] int * histogram; // empty array of at least 5 elements unsigned int levels; // e.g., 6 (for 5 bins) float lower_level; // e.g., 0.0 float upper_level; // e.g., 10.0 size_t temporary_storage_size_bytes; void * temporary_storage_ptr = nullptr; // Get required size of the temporary storage rocprim::histogram_even( temporary_storage_ptr, temporary_storage_size_bytes, samples, columns, rows, row_stride_bytes, histogram, levels, lower_level, upper_level ); // allocate temporary storage hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes); // compute histogram rocprim::histogram_even( temporary_storage_ptr, temporary_storage_size_bytes, samples, columns, rows, row_stride_bytes, histogram, levels, lower_level, upper_level ); // histogram: [3, 0, 1, 0, 2]
- Template Parameters:
Config – - [optional] configuration of the primitive. It has to be
histogram_config
or a class derived from it.SampleIterator – - random-access iterator type of the input range. Must meet the requirements of a C++ InputIterator concept. It can be a simple pointer type.
Counter – - integer type for histogram bin counters.
Level – - type of histogram boundaries (levels)
- Parameters:
temporary_storage – [in] - pointer to a device-accessible temporary storage. When a null pointer is passed, the required allocation size (in bytes) is written to
storage_size
and function returns without performing the reduction operation.storage_size – [inout] - reference to a size (in bytes) of
temporary_storage
.samples – [in] - iterator to the first element in the range of input samples.
columns – [in] - number of elements in each row of the region.
rows – [in] - number of rows of the region.
row_stride_bytes – [in] - number of bytes between starts of consecutive rows of the region.
histogram – [out] - pointer to the first element in the histogram range.
levels – [in] - number of boundaries (levels) for histogram bins.
lower_level – [in] - lower sample value bound (inclusive) for the first histogram bin.
upper_level – [in] - upper sample value bound (exclusive) for the last histogram bin.
stream – [in] - [optional] HIP stream object. Default is
0
(default stream).debug_synchronous – [in] - [optional] If true, synchronization after every kernel launch is forced in order to check for errors. Default value is
false
.
- Returns:
hipSuccess
(0
) after successful histogram operation; otherwise a HIP runtime error of typehipError_t
.
multi_histogram_even#
-
template<unsigned int Channels, unsigned int ActiveChannels, class Config = default_config, class SampleIterator, class Counter, class Level>
inline hipError_t rocprim::multi_histogram_even(void *temporary_storage, size_t &storage_size, SampleIterator samples, unsigned int size, Counter *histogram[ActiveChannels], unsigned int levels[ActiveChannels], Level lower_level[ActiveChannels], Level upper_level[ActiveChannels], hipStream_t stream = 0, bool debug_synchronous = false)# Computes histograms from a sequence of multi-channel samples using equal-width bins.
The input is a sequence of pixel structures, where each pixel comprises a record of
Channels
consecutive data samples (e.g.,Channels
= 4 for RGBA samples).The first
ActiveChannels
channels of totalChannels
channels will be used for computing histograms (e.g.,ActiveChannels
= 3 for computing histograms of only RGB from RGBA samples).For channel the number of histogram bins is (
levels
[i] - 1).For channel bins are evenly-segmented and include the same width of sample values: (
upper_level
[i] -lower_level
[i]) / (levels
[i] - 1).Returns the required size of
temporary_storage
instorage_size
iftemporary_storage
in a null pointer.
- Example
In this example histograms for 3 channels (RGB) are computed on an array of 8-bit RGBA samples.
#include <rocprim/rocprim.hpp> // Prepare input and output (declare pointers, allocate device memory etc.) unsigned int size; // e.g., 8 unsigned char * samples; // e.g., [(3, 1, 5, 255), (3, 1, 5, 255), (4, 2, 6, 127), (3, 2, 6, 127), // (0, 0, 0, 100), (0, 1, 0, 100), (0, 0, 1, 255), (0, 1, 1, 255)] int * histogram[3]; // 3 empty arrays of at least 256 elements each unsigned int levels[3]; // e.g., [257, 257, 257] (for 256 bins) int lower_level[3]; // e.g., [0, 0, 0] int upper_level[3]; // e.g., [256, 256, 256] size_t temporary_storage_size_bytes; void * temporary_storage_ptr = nullptr; // Get required size of the temporary storage rocprim::multi_histogram_even<4, 3>( temporary_storage_ptr, temporary_storage_size_bytes, samples, size, histogram, levels, lower_level, upper_level ); // allocate temporary storage hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes); // compute histograms rocprim::multi_histogram_even<4, 3>( temporary_storage_ptr, temporary_storage_size_bytes, samples, size, histogram, levels, lower_level, upper_level ); // histogram: [[4, 0, 0, 3, 1, 0, 0, ..., 0], // [2, 4, 2, 0, 0, 0, 0, ..., 0], // [2, 2, 0, 0, 0, 2, 2, ..., 0]]
- Template Parameters:
Channels – - number of channels interleaved in the input samples.
ActiveChannels – - number of channels being used for computing histograms.
Config – - [optional] configuration of the primitive. It has to be
histogram_config
or a class derived from it.SampleIterator – - random-access iterator type of the input range. Must meet the requirements of a C++ InputIterator concept. It can be a simple pointer type.
Counter – - integer type for histogram bin counters.
Level – - type of histogram boundaries (levels)
- Parameters:
temporary_storage – [in] - pointer to a device-accessible temporary storage. When a null pointer is passed, the required allocation size (in bytes) is written to
storage_size
and function returns without performing the reduction operation.storage_size – [inout] - reference to a size (in bytes) of
temporary_storage
.samples – [in] - iterator to the first element in the range of input samples.
size – [in] - number of pixels in the samples range.
histogram – [out] - pointers to the first element in the histogram range, one for each active channel.
levels – [in] - number of boundaries (levels) for histogram bins in each active channel.
lower_level – [in] - lower sample value bound (inclusive) for the first histogram bin in each active channel.
upper_level – [in] - upper sample value bound (exclusive) for the last histogram bin in each active channel.
stream – [in] - [optional] HIP stream object. Default is
0
(default stream).debug_synchronous – [in] - [optional] If true, synchronization after every kernel launch is forced in order to check for errors. Default value is
false
.
- Returns:
hipSuccess
(0
) after successful histogram operation; otherwise a HIP runtime error of typehipError_t
.
-
template<unsigned int Channels, unsigned int ActiveChannels, class Config = default_config, class SampleIterator, class Counter, class Level>
inline hipError_t rocprim::multi_histogram_even(void *temporary_storage, size_t &storage_size, SampleIterator samples, unsigned int columns, unsigned int rows, size_t row_stride_bytes, Counter *histogram[ActiveChannels], unsigned int levels[ActiveChannels], Level lower_level[ActiveChannels], Level upper_level[ActiveChannels], hipStream_t stream = 0, bool debug_synchronous = false)# Computes histograms from a two-dimensional region of multi-channel samples using equal-width bins.
The two-dimensional region of interest within
samples
can be specified using thecolumns
,rows
androw_stride_bytes
parameters.The row stride must be a whole multiple of the sample data type size, i.e.,
(row_stride_bytes % sizeof(std::iterator_traits<SampleIterator>::value_type)) == 0
.The input is a sequence of pixel structures, where each pixel comprises a record of
Channels
consecutive data samples (e.g.,Channels
= 4 for RGBA samples).The first
ActiveChannels
channels of totalChannels
channels will be used for computing histograms (e.g.,ActiveChannels
= 3 for computing histograms of only RGB from RGBA samples).For channel the number of histogram bins is (
levels
[i] - 1).For channel bins are evenly-segmented and include the same width of sample values: (
upper_level
[i] -lower_level
[i]) / (levels
[i] - 1).Returns the required size of
temporary_storage
instorage_size
iftemporary_storage
in a null pointer.
- Example
In this example histograms for 3 channels (RGB) are computed on an array of 8-bit RGBA samples.
#include <rocprim/rocprim.hpp> // Prepare input and output (declare pointers, allocate device memory etc.) unsigned int columns; // e.g., 4 unsigned int rows; // e.g., 2 size_t row_stride_bytes; // e.g., 5 * sizeof(unsigned char) unsigned char * samples; // e.g., [(3, 1, 5, 255), (3, 1, 5, 255), (4, 2, 6, 127), (3, 2, 6, 127), (-, -, -, -), // (0, 0, 0, 100), (0, 1, 0, 100), (0, 0, 1, 255), (0, 1, 1, 255), (-, -, -, -)] int * histogram[3]; // 3 empty arrays of at least 256 elements each unsigned int levels[3]; // e.g., [257, 257, 257] (for 256 bins) int lower_level[3]; // e.g., [0, 0, 0] int upper_level[3]; // e.g., [256, 256, 256] size_t temporary_storage_size_bytes; void * temporary_storage_ptr = nullptr; // Get required size of the temporary storage rocprim::multi_histogram_even<4, 3>( temporary_storage_ptr, temporary_storage_size_bytes, samples, columns, rows, row_stride_bytes, histogram, levels, lower_level, upper_level ); // allocate temporary storage hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes); // compute histograms rocprim::multi_histogram_even<4, 3>( temporary_storage_ptr, temporary_storage_size_bytes, samples, columns, rows, row_stride_bytes, histogram, levels, lower_level, upper_level ); // histogram: [[4, 0, 0, 3, 1, 0, 0, ..., 0], // [2, 4, 2, 0, 0, 0, 0, ..., 0], // [2, 2, 0, 0, 0, 2, 2, ..., 0]]
- Template Parameters:
Channels – - number of channels interleaved in the input samples.
ActiveChannels – - number of channels being used for computing histograms.
Config – - [optional] configuration of the primitive. It has to be
histogram_config
or a class derived from it.SampleIterator – - random-access iterator type of the input range. Must meet the requirements of a C++ InputIterator concept. It can be a simple pointer type.
Counter – - integer type for histogram bin counters.
Level – - type of histogram boundaries (levels)
- Parameters:
temporary_storage – [in] - pointer to a device-accessible temporary storage. When a null pointer is passed, the required allocation size (in bytes) is written to
storage_size
and function returns without performing the reduction operation.storage_size – [inout] - reference to a size (in bytes) of
temporary_storage
.samples – [in] - iterator to the first element in the range of input samples.
columns – [in] - number of elements in each row of the region.
rows – [in] - number of rows of the region.
row_stride_bytes – [in] - number of bytes between starts of consecutive rows of the region.
histogram – [out] - pointers to the first element in the histogram range, one for each active channel.
levels – [in] - number of boundaries (levels) for histogram bins in each active channel.
lower_level – [in] - lower sample value bound (inclusive) for the first histogram bin in each active channel.
upper_level – [in] - upper sample value bound (exclusive) for the last histogram bin in each active channel.
stream – [in] - [optional] HIP stream object. Default is
0
(default stream).debug_synchronous – [in] - [optional] If true, synchronization after every kernel launch is forced in order to check for errors. Default value is
false
.
- Returns:
hipSuccess
(0
) after successful histogram operation; otherwise a HIP runtime error of typehipError_t
.
histogram_range#
-
template<class Config = default_config, class SampleIterator, class Counter, class Level>
inline hipError_t rocprim::histogram_range(void *temporary_storage, size_t &storage_size, SampleIterator samples, unsigned int size, Counter *histogram, unsigned int levels, Level *level_values, hipStream_t stream = 0, bool debug_synchronous = false)# Computes a histogram from a sequence of samples using the specified bin boundary levels.
The number of histogram bins is (
levels
- 1).The range for bin is [
level_values[j]
,level_values[j+1]
).Returns the required size of
temporary_storage
instorage_size
iftemporary_storage
in a null pointer.
- Example
In this example a device-level histogram of 5 bins is computed on an array of float samples.
#include <rocprim/rocprim.hpp> // Prepare input and output (declare pointers, allocate device memory etc.) unsigned int size; // e.g., 8 float * samples; // e.g., [-10.0, 0.3, 9.5, 8.1, 1.5, 1.9, 100.0, 5.1] int * histogram; // empty array of at least 5 elements unsigned int levels; // e.g., 6 (for 5 bins) float * level_values; // e.g., [0.0, 1.0, 5.0, 10.0, 20.0, 50.0] size_t temporary_storage_size_bytes; void * temporary_storage_ptr = nullptr; // Get required size of the temporary storage rocprim::histogram_range( temporary_storage_ptr, temporary_storage_size_bytes, samples, size, histogram, levels, level_values ); // allocate temporary storage hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes); // compute histogram rocprim::histogram_range( temporary_storage_ptr, temporary_storage_size_bytes, samples, size, histogram, levels, level_values ); // histogram: [1, 2, 3, 0, 0]
- Template Parameters:
Config – - [optional] configuration of the primitive. It has to be
histogram_config
or a class derived from it.SampleIterator – - random-access iterator type of the input range. Must meet the requirements of a C++ InputIterator concept. It can be a simple pointer type.
Counter – - integer type for histogram bin counters.
Level – - type of histogram boundaries (levels)
- Parameters:
temporary_storage – [in] - pointer to a device-accessible temporary storage. When a null pointer is passed, the required allocation size (in bytes) is written to
storage_size
and function returns without performing the reduction operation.storage_size – [inout] - reference to a size (in bytes) of
temporary_storage
.samples – [in] - iterator to the first element in the range of input samples.
size – [in] - number of elements in the samples range.
histogram – [out] - pointer to the first element in the histogram range.
levels – [in] - number of boundaries (levels) for histogram bins.
level_values – [in] - pointer to the array of bin boundaries.
stream – [in] - [optional] HIP stream object. Default is
0
(default stream).debug_synchronous – [in] - [optional] If true, synchronization after every kernel launch is forced in order to check for errors. Default value is
false
.
- Returns:
hipSuccess
(0
) after successful histogram operation; otherwise a HIP runtime error of typehipError_t
.
-
template<class Config = default_config, class SampleIterator, class Counter, class Level>
inline hipError_t rocprim::histogram_range(void *temporary_storage, size_t &storage_size, SampleIterator samples, unsigned int columns, unsigned int rows, size_t row_stride_bytes, Counter *histogram, unsigned int levels, Level *level_values, hipStream_t stream = 0, bool debug_synchronous = false)# Computes a histogram from a two-dimensional region of samples using the specified bin boundary levels.
The two-dimensional region of interest within
samples
can be specified using thecolumns
,rows
androw_stride_bytes
parameters.The row stride must be a whole multiple of the sample data type size, i.e.,
(row_stride_bytes % sizeof(std::iterator_traits<SampleIterator>::value_type)) == 0
.The number of histogram bins is (
levels
- 1).The range for bin is [
level_values[j]
,level_values[j+1]
).Returns the required size of
temporary_storage
instorage_size
iftemporary_storage
in a null pointer.
- Example
In this example a device-level histogram of 5 bins is computed on an array of float samples.
#include <rocprim/rocprim.hpp> // Prepare input and output (declare pointers, allocate device memory etc.) unsigned int columns; // e.g., 4 unsigned int rows; // e.g., 2 size_t row_stride_bytes; // e.g., 6 * sizeof(float) float * samples; // e.g., [-10.0, 0.3, 9.5, 8.1, 1.5, 1.9, 100.0, 5.1] int * histogram; // empty array of at least 5 elements unsigned int levels; // e.g., 6 (for 5 bins) float level_values; // e.g., [0.0, 1.0, 5.0, 10.0, 20.0, 50.0] size_t temporary_storage_size_bytes; void * temporary_storage_ptr = nullptr; // Get required size of the temporary storage rocprim::histogram_range( temporary_storage_ptr, temporary_storage_size_bytes, samples, columns, rows, row_stride_bytes, histogram, levels, level_values ); // allocate temporary storage hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes); // compute histogram rocprim::histogram_range( temporary_storage_ptr, temporary_storage_size_bytes, samples, columns, rows, row_stride_bytes, histogram, levels, level_values ); // histogram: [1, 2, 3, 0, 0]
- Template Parameters:
Config – - [optional] configuration of the primitive. It has to be
histogram_config
or a class derived from it.SampleIterator – - random-access iterator type of the input range. Must meet the requirements of a C++ InputIterator concept. It can be a simple pointer type.
Counter – - integer type for histogram bin counters.
Level – - type of histogram boundaries (levels)
- Parameters:
temporary_storage – [in] - pointer to a device-accessible temporary storage. When a null pointer is passed, the required allocation size (in bytes) is written to
storage_size
and function returns without performing the reduction operation.storage_size – [inout] - reference to a size (in bytes) of
temporary_storage
.samples – [in] - iterator to the first element in the range of input samples.
columns – [in] - number of elements in each row of the region.
rows – [in] - number of rows of the region.
row_stride_bytes – [in] - number of bytes between starts of consecutive rows of the region.
histogram – [out] - pointer to the first element in the histogram range.
levels – [in] - number of boundaries (levels) for histogram bins.
level_values – [in] - pointer to the array of bin boundaries.
stream – [in] - [optional] HIP stream object. Default is
0
(default stream).debug_synchronous – [in] - [optional] If true, synchronization after every kernel launch is forced in order to check for errors. Default value is
false
.
- Returns:
hipSuccess
(0
) after successful histogram operation; otherwise a HIP runtime error of typehipError_t
.
multi_histogram_range#
-
template<unsigned int Channels, unsigned int ActiveChannels, class Config = default_config, class SampleIterator, class Counter, class Level>
inline hipError_t rocprim::multi_histogram_range(void *temporary_storage, size_t &storage_size, SampleIterator samples, unsigned int size, Counter *histogram[ActiveChannels], unsigned int levels[ActiveChannels], Level *level_values[ActiveChannels], hipStream_t stream = 0, bool debug_synchronous = false)# Computes histograms from a sequence of multi-channel samples using the specified bin boundary levels.
The input is a sequence of pixel structures, where each pixel comprises a record of
Channels
consecutive data samples (e.g.,Channels
= 4 for RGBA samples).The first
ActiveChannels
channels of totalChannels
channels will be used for computing histograms (e.g.,ActiveChannels
= 3 for computing histograms of only RGB from RGBA samples).For channel the number of histogram bins is (
levels
[i] - 1).For channel the range for bin is [
level_values[i][j]
,level_values[i][j+1]
).Returns the required size of
temporary_storage
instorage_size
iftemporary_storage
in a null pointer.
- Example
In this example histograms for 3 channels (RGB) are computed on an array of 8-bit RGBA samples.
#include <rocprim/rocprim.hpp> // Prepare input and output (declare pointers, allocate device memory etc.) unsigned int size; // e.g., 8 unsigned char * samples; // e.g., [(0, 0, 80, 255), (120, 0, 80, 255), (123, 0, 82, 127), (10, 1, 83, 127), // (51, 1, 8, 100), (52, 1, 8, 100), (53, 0, 81, 255), (54, 50, 81, 255)] int * histogram[3]; // 3 empty arrays of at least 256 elements each unsigned int levels[3]; // e.g., [4, 4, 3] int * level_values[3]; // e.g., [[0, 50, 100, 200], [0, 20, 40, 60], [0, 10, 100]] size_t temporary_storage_size_bytes; void * temporary_storage_ptr = nullptr; // Get required size of the temporary storage rocprim::multi_histogram_range<4, 3>( temporary_storage_ptr, temporary_storage_size_bytes, samples, size, histogram, levels, level_values ); // allocate temporary storage hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes); // compute histograms rocprim::multi_histogram_range<4, 3>( temporary_storage_ptr, temporary_storage_size_bytes, samples, size, histogram, levels, level_values ); // histogram: [[2, 4, 2], [7, 0, 1], [2, 6]]
- Template Parameters:
Channels – - number of channels interleaved in the input samples.
ActiveChannels – - number of channels being used for computing histograms.
Config – - [optional] configuration of the primitive. It has to be
histogram_config
or a class derived from it.SampleIterator – - random-access iterator type of the input range. Must meet the requirements of a C++ InputIterator concept. It can be a simple pointer type.
Counter – - integer type for histogram bin counters.
Level – - type of histogram boundaries (levels)
- Parameters:
temporary_storage – [in] - pointer to a device-accessible temporary storage. When a null pointer is passed, the required allocation size (in bytes) is written to
storage_size
and function returns without performing the reduction operation.storage_size – [inout] - reference to a size (in bytes) of
temporary_storage
.samples – [in] - iterator to the first element in the range of input samples.
size – [in] - number of pixels in the samples range.
histogram – [out] - pointers to the first element in the histogram range, one for each active channel.
levels – [in] - number of boundaries (levels) for histogram bins in each active channel.
level_values – [in] - pointer to the array of bin boundaries for each active channel.
stream – [in] - [optional] HIP stream object. Default is
0
(default stream).debug_synchronous – [in] - [optional] If true, synchronization after every kernel launch is forced in order to check for errors. Default value is
false
.
- Returns:
hipSuccess
(0
) after successful histogram operation; otherwise a HIP runtime error of typehipError_t
.
-
template<unsigned int Channels, unsigned int ActiveChannels, class Config = default_config, class SampleIterator, class Counter, class Level>
inline hipError_t rocprim::multi_histogram_range(void *temporary_storage, size_t &storage_size, SampleIterator samples, unsigned int columns, unsigned int rows, size_t row_stride_bytes, Counter *histogram[ActiveChannels], unsigned int levels[ActiveChannels], Level *level_values[ActiveChannels], hipStream_t stream = 0, bool debug_synchronous = false)# Computes histograms from a two-dimensional region of multi-channel samples using the specified bin boundary levels.
The two-dimensional region of interest within
samples
can be specified using thecolumns
,rows
androw_stride_bytes
parameters.The row stride must be a whole multiple of the sample data type size, i.e.,
(row_stride_bytes % sizeof(std::iterator_traits<SampleIterator>::value_type)) == 0
.The input is a sequence of pixel structures, where each pixel comprises a record of
Channels
consecutive data samples (e.g.,Channels
= 4 for RGBA samples).The first
ActiveChannels
channels of totalChannels
channels will be used for computing histograms (e.g.,ActiveChannels
= 3 for computing histograms of only RGB from RGBA samples).For channel the number of histogram bins is (
levels
[i] - 1).For channel the range for bin is [
level_values[i][j]
,level_values[i][j+1]
).Returns the required size of
temporary_storage
instorage_size
iftemporary_storage
in a null pointer.
- Example
In this example histograms for 3 channels (RGB) are computed on an array of 8-bit RGBA samples.
#include <rocprim/rocprim.hpp> // Prepare input and output (declare pointers, allocate device memory etc.) unsigned int columns; // e.g., 4 unsigned int rows; // e.g., 2 size_t row_stride_bytes; // e.g., 5 * sizeof(unsigned char) unsigned char * samples; // e.g., [(0, 0, 80, 0), (120, 0, 80, 0), (123, 0, 82, 0), (10, 1, 83, 0), (-, -, -, -), // (51, 1, 8, 0), (52, 1, 8, 0), (53, 0, 81, 0), (54, 50, 81, 0), (-, -, -, -)] int * histogram[3]; // 3 empty arrays unsigned int levels[3]; // e.g., [4, 4, 3] int * level_values[3]; // e.g., [[0, 50, 100, 200], [0, 20, 40, 60], [0, 10, 100]] size_t temporary_storage_size_bytes; void * temporary_storage_ptr = nullptr; // Get required size of the temporary storage rocprim::multi_histogram_range<4, 3>( temporary_storage_ptr, temporary_storage_size_bytes, samples, columns, rows, row_stride_bytes, histogram, levels, level_values ); // allocate temporary storage hipMalloc(&temporary_storage_ptr, temporary_storage_size_bytes); // compute histograms rocprim::multi_histogram_range<4, 3>( temporary_storage_ptr, temporary_storage_size_bytes, samples, columns, rows, row_stride_bytes, histogram, levels, level_values ); // histogram: [[2, 4, 2], [7, 0, 1], [2, 6]]
- Template Parameters:
Channels – - number of channels interleaved in the input samples.
ActiveChannels – - number of channels being used for computing histograms.
Config – - [optional] configuration of the primitive. It has to be
histogram_config
or a class derived from it.SampleIterator – - random-access iterator type of the input range. Must meet the requirements of a C++ InputIterator concept. It can be a simple pointer type.
Counter – - integer type for histogram bin counters.
Level – - type of histogram boundaries (levels)
- Parameters:
temporary_storage – [in] - pointer to a device-accessible temporary storage. When a null pointer is passed, the required allocation size (in bytes) is written to
storage_size
and function returns without performing the reduction operation.storage_size – [inout] - reference to a size (in bytes) of
temporary_storage
.samples – [in] - iterator to the first element in the range of input samples.
columns – [in] - number of elements in each row of the region.
rows – [in] - number of rows of the region.
row_stride_bytes – [in] - number of bytes between starts of consecutive rows of the region.
histogram – [out] - pointers to the first element in the histogram range, one for each active channel.
levels – [in] - number of boundaries (levels) for histogram bins in each active channel.
level_values – [in] - pointer to the array of bin boundaries for each active channel.
stream – [in] - [optional] HIP stream object. Default is
0
(default stream).debug_synchronous – [in] - [optional] If true, synchronization after every kernel launch is forced in order to check for errors. Default value is
false
.
- Returns:
hipSuccess
(0
) after successful histogram operation; otherwise a HIP runtime error of typehipError_t
.