Histogram

Histogram#

Configuring the kernel#

template<class HistogramConfig, unsigned int MaxGridSize = 1024, unsigned int SharedImplMaxBins = 2048, unsigned int SharedImplHistograms = 3> struct histogram_config#

Configuration of device-level histogram operation.

Template Parameters:

HistogramConfig – - configuration of histogram kernel. Must be kernel_config.
MaxGridSize – - maximim 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.

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. It can be histogram_config or a custom class with the same members.
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. It can be histogram_config or a custom class with the same members.
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.

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 can be histogram_config or a custom class with the same members.
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.

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 can be histogram_config or a custom class with the same members.
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. It can be histogram_config or a custom class with the same members.
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. It can be histogram_config or a custom class with the same members.
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.

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 can be histogram_config or a custom class with the same members.
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.

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 can be histogram_config or a custom class with the same members.
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#