Histogram

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 in storage_size if temporary_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, must be default_config or histogram_config.
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 type hipError_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 the columns, rows and row_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 in storage_size if temporary_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, must be default_config or histogram_config.
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 type hipError_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 total Channels 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 in storage_size if temporary_storage in a null pointer.

Notes

Currently the Channels template parameter has no strict restriction on its value. However, internally a vector type of elements of type SampleIterator and length Channels is used to represent the input items, so the amount of local memory available will limit the range of possible values for this template parameter.
ActiveChannels must be less or equal than Channels.

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, must be default_config or histogram_config.
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 type hipError_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 the columns, rows and row_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 total Channels 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 in storage_size if temporary_storage in a null pointer.

Notes

Currently the Channels template parameter has no strict restriction on its value. However, internally a vector type of elements of type SampleIterator and length Channels is used to represent the input items, so the amount of local memory available will limit the range of possible values for this template parameter.
ActiveChannels must be less or equal than Channels.

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, must be default_config or histogram_config.
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 type hipError_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 in storage_size if temporary_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, must be default_config or histogram_config.
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 type hipError_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 the columns, rows and row_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 in storage_size if temporary_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, must be default_config or histogram_config.
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 type hipError_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 total Channels 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 in storage_size if temporary_storage in a null pointer.

Notes

Currently the Channels template parameter has no strict restriction on its value. However, internally a vector type of elements of type SampleIterator and length Channels is used to represent the input items, so the amount of local memory available will limit the range of possible values for this template parameter.
ActiveChannels must be less or equal than Channels.

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, must be default_config or histogram_config.
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 type hipError_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 the columns, rows and row_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 total Channels 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 in storage_size if temporary_storage in a null pointer.

Notes

Currently the Channels template parameter has no strict restriction on its value. However, internally a vector type of elements of type SampleIterator and length Channels is used to represent the input items, so the amount of local memory available will limit the range of possible values for this template parameter.
ActiveChannels must be less or equal than Channels.

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, must be default_config or histogram_config.
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 type hipError_t.

Histogram

Contents

Histogram#

Configuring the kernel#

histogram_even#

multi_histogram_even#

histogram_range#

multi_histogram_range#