Detailed description

DeviceHistogram provides device-wide parallel operations for constructing histogram(s) from a sequence of samples data residing within device-accessible memory.

.

Overview: A histogram counts the number of observations that fall into each of the disjoint categories (known as bins).

Usage Considerations

Dynamic parallelism. DeviceHistogram methods can be called within kernel code on devices in which CUDA dynamic parallelism is supported.

Static Public Methods
Evenly-segmented bin ranges
template<typename SampleIteratorT , typename CounterT , typename LevelT , typename OffsetT >
static CUB_RUNTIME_FUNCTION cudaError_t	HistogramEven (void d_temp_storage, size_t &temp_storage_bytes, SampleIteratorT d_samples, CounterT d_histogram, int num_levels, LevelT lower_level, LevelT upper_level, OffsetT num_samples, cudaStream_t stream=0, bool debug_synchronous=false)
	Computes an intensity histogram from a sequence of data samples using equal-width bins. More...

template<typename SampleIteratorT , typename CounterT , typename LevelT , typename OffsetT >
static CUB_RUNTIME_FUNCTION cudaError_t	HistogramEven (void d_temp_storage, size_t &temp_storage_bytes, SampleIteratorT d_samples, CounterT d_histogram, int num_levels, LevelT lower_level, LevelT upper_level, OffsetT num_row_samples, OffsetT num_rows, size_t row_stride_bytes, cudaStream_t stream=0, bool debug_synchronous=false)
	Computes an intensity histogram from a sequence of data samples using equal-width bins. More...

template<int NUM_CHANNELS, int NUM_ACTIVE_CHANNELS, typename SampleIteratorT , typename CounterT , typename LevelT , typename OffsetT >
static CUB_RUNTIME_FUNCTION cudaError_t	MultiHistogramEven (void d_temp_storage, size_t &temp_storage_bytes, SampleIteratorT d_samples, CounterT d_histogram[NUM_ACTIVE_CHANNELS], int num_levels[NUM_ACTIVE_CHANNELS], LevelT lower_level[NUM_ACTIVE_CHANNELS], LevelT upper_level[NUM_ACTIVE_CHANNELS], OffsetT num_pixels, cudaStream_t stream=0, bool debug_synchronous=false)
	Computes per-channel intensity histograms from a sequence of multi-channel "pixel" data samples using equal-width bins. More...

template<int NUM_CHANNELS, int NUM_ACTIVE_CHANNELS, typename SampleIteratorT , typename CounterT , typename LevelT , typename OffsetT >
static CUB_RUNTIME_FUNCTION cudaError_t	MultiHistogramEven (void d_temp_storage, size_t &temp_storage_bytes, SampleIteratorT d_samples, CounterT d_histogram[NUM_ACTIVE_CHANNELS], int num_levels[NUM_ACTIVE_CHANNELS], LevelT lower_level[NUM_ACTIVE_CHANNELS], LevelT upper_level[NUM_ACTIVE_CHANNELS], OffsetT num_row_pixels, OffsetT num_rows, size_t row_stride_bytes, cudaStream_t stream=0, bool debug_synchronous=false)
	Computes per-channel intensity histograms from a sequence of multi-channel "pixel" data samples using equal-width bins. More...

Custom bin ranges
template<typename SampleIteratorT , typename CounterT , typename LevelT , typename OffsetT >
static CUB_RUNTIME_FUNCTION cudaError_t	HistogramRange (void d_temp_storage, size_t &temp_storage_bytes, SampleIteratorT d_samples, CounterT d_histogram, int num_levels, LevelT *d_levels, OffsetT num_samples, cudaStream_t stream=0, bool debug_synchronous=false)
	Computes an intensity histogram from a sequence of data samples using the specified bin boundary levels. More...

template<typename SampleIteratorT , typename CounterT , typename LevelT , typename OffsetT >
static CUB_RUNTIME_FUNCTION cudaError_t	HistogramRange (void d_temp_storage, size_t &temp_storage_bytes, SampleIteratorT d_samples, CounterT d_histogram, int num_levels, LevelT *d_levels, OffsetT num_row_samples, OffsetT num_rows, size_t row_stride_bytes, cudaStream_t stream=0, bool debug_synchronous=false)
	Computes an intensity histogram from a sequence of data samples using the specified bin boundary levels. More...

template<int NUM_CHANNELS, int NUM_ACTIVE_CHANNELS, typename SampleIteratorT , typename CounterT , typename LevelT , typename OffsetT >
static CUB_RUNTIME_FUNCTION cudaError_t	MultiHistogramRange (void d_temp_storage, size_t &temp_storage_bytes, SampleIteratorT d_samples, CounterT d_histogram[NUM_ACTIVE_CHANNELS], int num_levels[NUM_ACTIVE_CHANNELS], LevelT *d_levels[NUM_ACTIVE_CHANNELS], OffsetT num_pixels, cudaStream_t stream=0, bool debug_synchronous=false)
	Computes per-channel intensity histograms from a sequence of multi-channel "pixel" data samples using the specified bin boundary levels. More...

template<int NUM_CHANNELS, int NUM_ACTIVE_CHANNELS, typename SampleIteratorT , typename CounterT , typename LevelT , typename OffsetT >
static CUB_RUNTIME_FUNCTION cudaError_t	MultiHistogramRange (void d_temp_storage, size_t &temp_storage_bytes, SampleIteratorT d_samples, CounterT d_histogram[NUM_ACTIVE_CHANNELS], int num_levels[NUM_ACTIVE_CHANNELS], LevelT *d_levels[NUM_ACTIVE_CHANNELS], OffsetT num_row_pixels, OffsetT num_rows, size_t row_stride_bytes, cudaStream_t stream=0, bool debug_synchronous=false)
	Computes per-channel intensity histograms from a sequence of multi-channel "pixel" data samples using the specified bin boundary levels. More...

Member Function Documentation

template<typename SampleIteratorT , typename CounterT , typename LevelT , typename OffsetT >

static CUB_RUNTIME_FUNCTION cudaError_t cub::DeviceHistogram::HistogramEven	(	void *	d_temp_storage,
		size_t &	temp_storage_bytes,
		SampleIteratorT	d_samples,
		CounterT *	d_histogram,
		int	num_levels,
		LevelT	lower_level,
		LevelT	upper_level,
		OffsetT	num_samples,
		cudaStream_t	stream = `0`,
		bool	debug_synchronous = `false`
	)

inlinestatic

Computes an intensity histogram from a sequence of data samples using equal-width bins.

The number of histogram bins is (num_levels - 1)
All bins comprise the same width of sample values: (upper_level - lower_level) / (num_levels - 1)
When d_temp_storage is NULL, no work is done and the required allocation size is returned in temp_storage_bytes.

Snippet: The code snippet below illustrates the computation of a six-bin histogram from a sequence of float samples

: #include <cub/cub.cuh> // or equivalently <cub/device/device_histogram.cuh>

// Declare, allocate, and initialize device-accessible pointers for input samples and

// output histogram

int num_samples; // e.g., 10

float* d_samples; // e.g., [2.2, 6.1, 7.1, 2.9, 3.5, 0.3, 2.9, 2.1, 6.1, 999.5]

int* d_histogram; // e.g., [ -, -, -, -, -, -]

int num_levels; // e.g., 7 (seven level boundaries for six bins)

float lower_level; // e.g., 0.0 (lower sample value boundary of lowest bin)

float upper_level; // e.g., 12.0 (upper sample value boundary of upper bin)

...

// Determine temporary device storage requirements

void* d_temp_storage = NULL;

size_t temp_storage_bytes = 0;

cub::DeviceHistogram::HistogramEven(d_temp_storage, temp_storage_bytes,

d_samples, d_histogram, num_levels, lower_level, upper_level, num_samples);

// Allocate temporary storage

cudaMalloc(&d_temp_storage, temp_storage_bytes);

// Compute histograms

cub::DeviceHistogram::HistogramEven(d_temp_storage, temp_storage_bytes,

d_samples, d_histogram, num_levels, lower_level, upper_level, num_samples);

// d_histogram <-- [1, 5, 0, 3, 0, 0];

Template Parameters

SampleIteratorT	[inferred] Random-access input iterator type for reading input samples. (may be a simple pointer type)
CounterT	[inferred] Integer type for histogram bin counters
LevelT	[inferred] Type for specifying boundaries (levels)
OffsetT	[inferred] Signed integer type for sequence offsets, list lengths, pointer differences, etc. (Consider using 32-bit values as offsets/lengths/etc. For example, `int` will typically yeild better performance than `size_t` in 64-bit memory mode.)

The sample value type of the input iterator

Parameters

[in]	d_temp_storage	Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to `temp_storage_bytes` and no work is done.
[in,out]	temp_storage_bytes	Reference to size in bytes of `d_temp_storage` allocation
[in]	d_samples	The pointer to the input sequence of data samples.
[out]	d_histogram	The pointer to the histogram counter output array of length `num_levels` - 1.
[in]	num_levels	The number of boundaries (levels) for delineating histogram samples. Implies that the number of bins is `num_levels` - 1.
[in]	lower_level	The lower sample value bound (inclusive) for the lowest histogram bin.
[in]	upper_level	The upper sample value bound (exclusive) for the highest histogram bin.
[in]	num_samples	The number of input samples (i.e., the length of `d_samples`)
[in]	stream	[optional] CUDA stream to launch kernels within. Default is stream₀.
[in]	debug_synchronous	[optional] Whether or not to synchronize the stream after every kernel launch to check for errors. May cause significant slowdown. Default is `false`.

template<typename SampleIteratorT , typename CounterT , typename LevelT , typename OffsetT >

static CUB_RUNTIME_FUNCTION cudaError_t cub::DeviceHistogram::HistogramEven	(	void *	d_temp_storage,
		size_t &	temp_storage_bytes,
		SampleIteratorT	d_samples,
		CounterT *	d_histogram,
		int	num_levels,
		LevelT	lower_level,
		LevelT	upper_level,
		OffsetT	num_row_samples,
		OffsetT	num_rows,
		size_t	row_stride_bytes,
		cudaStream_t	stream = `0`,
		bool	debug_synchronous = `false`
	)

inlinestatic

Computes an intensity histogram from a sequence of data samples using equal-width bins.

A two-dimensional region of interest within d_samples can be specified using the num_row_samples, num_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(SampleT)) == 0.
The number of histogram bins is (num_levels - 1)
All bins comprise the same width of sample values: (upper_level - lower_level) / (num_levels - 1)
When d_temp_storage is NULL, no work is done and the required allocation size is returned in temp_storage_bytes.

Snippet: The code snippet below illustrates the computation of a six-bin histogram from a 2x5 region of interest within a flattened 2x7 array of float samples.

: #include <cub/cub.cuh> // or equivalently <cub/device/device_histogram.cuh>

// Declare, allocate, and initialize device-accessible pointers for input samples and

// output histogram

int num_row_samples; // e.g., 5

int num_rows; // e.g., 2;

size_t row_stride_bytes; // e.g., 7 * sizeof(float)

float* d_samples; // e.g., [2.2, 6.1, 7.1, 2.9, 3.5, -, -,

// 0.3, 2.9, 2.1, 6.1, 999.5, -, -]

int* d_histogram; // e.g., [ -, -, -, -, -, -]

int num_levels; // e.g., 7 (seven level boundaries for six bins)

float lower_level; // e.g., 0.0 (lower sample value boundary of lowest bin)

float upper_level; // e.g., 12.0 (upper sample value boundary of upper bin)

...

// Determine temporary device storage requirements

void* d_temp_storage = NULL;

size_t temp_storage_bytes = 0;

cub::DeviceHistogram::HistogramEven(d_temp_storage, temp_storage_bytes,

d_samples, d_histogram, num_levels, lower_level, upper_level,

num_row_samples, num_rows, row_stride_bytes);

// Allocate temporary storage

cudaMalloc(&d_temp_storage, temp_storage_bytes);

// Compute histograms

cub::DeviceHistogram::HistogramEven(d_temp_storage, temp_storage_bytes, d_samples, d_histogram,

d_samples, d_histogram, num_levels, lower_level, upper_level,

num_row_samples, num_rows, row_stride_bytes);

// d_histogram <-- [1, 5, 0, 3, 0, 0];

Template Parameters

SampleIteratorT	[inferred] Random-access input iterator type for reading input samples. (may be a simple pointer type)
CounterT	[inferred] Integer type for histogram bin counters
LevelT	[inferred] Type for specifying boundaries (levels)
OffsetT	[inferred] Signed integer type for sequence offsets, list lengths, pointer differences, etc. (Consider using 32-bit values as offsets/lengths/etc. For example, `int` will typically yeild better performance than `size_t` in 64-bit memory mode.)

Parameters

[in]	d_temp_storage	Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to `temp_storage_bytes` and no work is done.
[in,out]	temp_storage_bytes	Reference to size in bytes of `d_temp_storage` allocation
[in]	d_samples	The pointer to the input sequence of data samples.
[out]	d_histogram	The pointer to the histogram counter output array of length `num_levels` - 1.
[in]	num_levels	The number of boundaries (levels) for delineating histogram samples. Implies that the number of bins is `num_levels` - 1.
[in]	lower_level	The lower sample value bound (inclusive) for the lowest histogram bin.
[in]	upper_level	The upper sample value bound (exclusive) for the highest histogram bin.
[in]	num_row_samples	The number of data samples per row in the region of interest
[in]	num_rows	The number of rows in the region of interest
[in]	row_stride_bytes	The number of bytes between starts of consecutive rows in the region of interest
[in]	stream	[optional] CUDA stream to launch kernels within. Default is stream₀.
[in]	debug_synchronous	[optional] Whether or not to synchronize the stream after every kernel launch to check for errors. May cause significant slowdown. Default is `false`.

template<int NUM_CHANNELS, int NUM_ACTIVE_CHANNELS, typename SampleIteratorT , typename CounterT , typename LevelT , typename OffsetT >

static CUB_RUNTIME_FUNCTION cudaError_t cub::DeviceHistogram::MultiHistogramEven	(	void *	d_temp_storage,
		size_t &	temp_storage_bytes,
		SampleIteratorT	d_samples,
		CounterT *	d_histogram[NUM_ACTIVE_CHANNELS],
		int	num_levels[NUM_ACTIVE_CHANNELS],
		LevelT	lower_level[NUM_ACTIVE_CHANNELS],
		LevelT	upper_level[NUM_ACTIVE_CHANNELS],
		OffsetT	num_pixels,
		cudaStream_t	stream = `0`,
		bool	debug_synchronous = `false`
	)

inlinestatic

Computes per-channel intensity histograms from a sequence of multi-channel "pixel" data samples using equal-width bins.

The input is a sequence of pixel structures, where each pixel comprises a record of NUM_CHANNELS consecutive data samples (e.g., an RGBA pixel).
Of the NUM_CHANNELS specified, the function will only compute histograms for the first NUM_ACTIVE_CHANNELS (e.g., only RGB histograms from RGBA pixel samples).
The number of histogram bins for channel_i is num_levels[i] - 1.
For channel_i, the range of values for all histogram bins have the same width: (upper_level[i] - lower_level[i]) / ( num_levels[i] - 1)
When d_temp_storage is NULL, no work is done and the required allocation size is returned in temp_storage_bytes.

Snippet: The code snippet below illustrates the computation of three 256-bin RGB histograms from a quad-channel sequence of RGBA pixels (8 bits per channel per pixel)

: #include <cub/cub.cuh> // or equivalently <cub/device/device_histogram.cuh>

// Declare, allocate, and initialize device-accessible pointers for input samples

// and output histograms

int num_pixels; // e.g., 5

unsigned char* d_samples; // e.g., [(2, 6, 7, 5), (3, 0, 2, 1), (7, 0, 6, 2),

// (0, 6, 7, 5), (3, 0, 2, 6)]

int* d_histogram[3]; // e.g., three device pointers to three device buffers,

// each allocated with 256 integer counters

int num_levels[3]; // e.g., {257, 257, 257};

unsigned int lower_level[3]; // e.g., {0, 0, 0};

unsigned int upper_level[3]; // e.g., {256, 256, 256};

...

// Determine temporary device storage requirements

void* d_temp_storage = NULL;

size_t temp_storage_bytes = 0;

cub::DeviceHistogram::MultiHistogramEven<4, 3>(d_temp_storage, temp_storage_bytes,

d_samples, d_histogram, num_levels, lower_level, upper_level, num_pixels);

// Allocate temporary storage

cudaMalloc(&d_temp_storage, temp_storage_bytes);

// Compute histograms

cub::DeviceHistogram::MultiHistogramEven<4, 3>(d_temp_storage, temp_storage_bytes,

d_samples, d_histogram, num_levels, lower_level, upper_level, num_pixels);

// d_histogram <-- [ [1, 0, 1, 2, 0, 0, 0, 1, 0, 0, 0, ..., 0],

// [0, 3, 0, 0, 0, 0, 2, 0, 0, 0, 0, ..., 0],

// [0, 0, 2, 0, 0, 0, 1, 2, 0, 0, 0, ..., 0] ]

Template Parameters

NUM_CHANNELS	Number of channels interleaved in the input data (may be greater than the number of channels being actively histogrammed)
NUM_ACTIVE_CHANNELS	[inferred] Number of channels actively being histogrammed
SampleIteratorT	[inferred] Random-access input iterator type for reading input samples. (may be a simple pointer type)
CounterT	[inferred] Integer type for histogram bin counters
LevelT	[inferred] Type for specifying boundaries (levels)
OffsetT	[inferred] Signed integer type for sequence offsets, list lengths, pointer differences, etc. (Consider using 32-bit values as offsets/lengths/etc. For example, `int` will typically yeild better performance than `size_t` in 64-bit memory mode.)

The sample value type of the input iterator

Parameters

[in]	d_temp_storage	Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to `temp_storage_bytes` and no work is done.
[in,out]	temp_storage_bytes	Reference to size in bytes of `d_temp_storage` allocation
[in]	d_samples	The pointer to the multi-channel input sequence of data samples. The samples from different channels are assumed to be interleaved (e.g., an array of 32-bit pixels where each pixel consists of four RGBA 8-bit samples).
[out]	d_histogram	The pointers to the histogram counter output arrays, one for each active channel. For channel_i, the allocation length of `d_histogram[i]` should be `num_levels[i]` - 1.
[in]	num_levels	The number of boundaries (levels) for delineating histogram samples in each active channel. Implies that the number of bins for channel_i is `num_levels[i]` - 1.
[in]	lower_level	The lower sample value bound (inclusive) for the lowest histogram bin in each active channel.
[in]	upper_level	The upper sample value bound (exclusive) for the highest histogram bin in each active channel.
[in]	num_pixels	The number of multi-channel pixels (i.e., the length of `d_samples` / NUM_CHANNELS)
[in]	stream	[optional] CUDA stream to launch kernels within. Default is stream₀.
[in]	debug_synchronous	[optional] Whether or not to synchronize the stream after every kernel launch to check for errors. May cause significant slowdown. Default is `false`.

template<int NUM_CHANNELS, int NUM_ACTIVE_CHANNELS, typename SampleIteratorT , typename CounterT , typename LevelT , typename OffsetT >

static CUB_RUNTIME_FUNCTION cudaError_t cub::DeviceHistogram::MultiHistogramEven	(	void *	d_temp_storage,
		size_t &	temp_storage_bytes,
		SampleIteratorT	d_samples,
		CounterT *	d_histogram[NUM_ACTIVE_CHANNELS],
		int	num_levels[NUM_ACTIVE_CHANNELS],
		LevelT	lower_level[NUM_ACTIVE_CHANNELS],
		LevelT	upper_level[NUM_ACTIVE_CHANNELS],
		OffsetT	num_row_pixels,
		OffsetT	num_rows,
		size_t	row_stride_bytes,
		cudaStream_t	stream = `0`,
		bool	debug_synchronous = `false`
	)

inlinestatic

Computes per-channel intensity histograms from a sequence of multi-channel "pixel" data samples using equal-width bins.

The input is a sequence of pixel structures, where each pixel comprises a record of NUM_CHANNELS consecutive data samples (e.g., an RGBA pixel).
Of the NUM_CHANNELS specified, the function will only compute histograms for the first NUM_ACTIVE_CHANNELS (e.g., only RGB histograms from RGBA pixel samples).
A two-dimensional region of interest within d_samples can be specified using the num_row_samples, num_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(SampleT)) == 0.
The number of histogram bins for channel_i is num_levels[i] - 1.
For channel_i, the range of values for all histogram bins have the same width: (upper_level[i] - lower_level[i]) / ( num_levels[i] - 1)
When d_temp_storage is NULL, no work is done and the required allocation size is returned in temp_storage_bytes.

Snippet: The code snippet below illustrates the computation of three 256-bin RGB histograms from a 2x3 region of interest of within a flattened 2x4 array of quad-channel RGBA pixels (8 bits per channel per pixel).

: #include <cub/cub.cuh> // or equivalently <cub/device/device_histogram.cuh>

// Declare, allocate, and initialize device-accessible pointers for input samples

// and output histograms

int num_row_pixels; // e.g., 3

int num_rows; // e.g., 2

size_t row_stride_bytes; // e.g., 4 * sizeof(unsigned char) * NUM_CHANNELS

unsigned char* d_samples; // e.g., [(2, 6, 7, 5), (3, 0, 2, 1), (7, 0, 6, 2), (-, -, -, -),

// (0, 6, 7, 5), (3, 0, 2, 6), (1, 1, 1, 1), (-, -, -, -)]

int* d_histogram[3]; // e.g., three device pointers to three device buffers,

// each allocated with 256 integer counters

int num_levels[3]; // e.g., {257, 257, 257};

unsigned int lower_level[3]; // e.g., {0, 0, 0};

unsigned int upper_level[3]; // e.g., {256, 256, 256};

...

// Determine temporary device storage requirements

void* d_temp_storage = NULL;

size_t temp_storage_bytes = 0;

cub::DeviceHistogram::MultiHistogramEven<4, 3>(d_temp_storage, temp_storage_bytes,

d_samples, d_histogram, num_levels, lower_level, upper_level,

num_row_pixels, num_rows, row_stride_bytes);

// Allocate temporary storage

cudaMalloc(&d_temp_storage, temp_storage_bytes);

// Compute histograms

cub::DeviceHistogram::MultiHistogramEven<4, 3>(d_temp_storage, temp_storage_bytes,

d_samples, d_histogram, num_levels, lower_level, upper_level,

num_row_pixels, num_rows, row_stride_bytes);

// d_histogram <-- [ [1, 1, 1, 2, 0, 0, 0, 1, 0, 0, 0, ..., 0],

// [0, 4, 0, 0, 0, 0, 2, 0, 0, 0, 0, ..., 0],

// [0, 1, 2, 0, 0, 0, 1, 2, 0, 0, 0, ..., 0] ]

Template Parameters

NUM_CHANNELS	Number of channels interleaved in the input data (may be greater than the number of channels being actively histogrammed)
NUM_ACTIVE_CHANNELS	[inferred] Number of channels actively being histogrammed
SampleIteratorT	[inferred] Random-access input iterator type for reading input samples. (may be a simple pointer type)
CounterT	[inferred] Integer type for histogram bin counters
LevelT	[inferred] Type for specifying boundaries (levels)
OffsetT	[inferred] Signed integer type for sequence offsets, list lengths, pointer differences, etc. (Consider using 32-bit values as offsets/lengths/etc. For example, `int` will typically yeild better performance than `size_t` in 64-bit memory mode.)

The sample value type of the input iterator

Parameters

[in]	d_temp_storage	Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to `temp_storage_bytes` and no work is done.
[in,out]	temp_storage_bytes	Reference to size in bytes of `d_temp_storage` allocation
[in]	d_samples	The pointer to the multi-channel input sequence of data samples. The samples from different channels are assumed to be interleaved (e.g., an array of 32-bit pixels where each pixel consists of four RGBA 8-bit samples).
[out]	d_histogram	The pointers to the histogram counter output arrays, one for each active channel. For channel_i, the allocation length of `d_histogram[i]` should be `num_levels[i]` - 1.
[in]	num_levels	The number of boundaries (levels) for delineating histogram samples in each active channel. Implies that the number of bins for channel_i is `num_levels[i]` - 1.
[in]	lower_level	The lower sample value bound (inclusive) for the lowest histogram bin in each active channel.
[in]	upper_level	The upper sample value bound (exclusive) for the highest histogram bin in each active channel.
[in]	num_row_pixels	The number of multi-channel pixels per row in the region of interest
[in]	num_rows	The number of rows in the region of interest
[in]	row_stride_bytes	The number of bytes between starts of consecutive rows in the region of interest
[in]	stream	[optional] CUDA stream to launch kernels within. Default is stream₀.
[in]	debug_synchronous	[optional] Whether or not to synchronize the stream after every kernel launch to check for errors. May cause significant slowdown. Default is `false`.

template<typename SampleIteratorT , typename CounterT , typename LevelT , typename OffsetT >

static CUB_RUNTIME_FUNCTION cudaError_t cub::DeviceHistogram::HistogramRange	(	void *	d_temp_storage,
		size_t &	temp_storage_bytes,
		SampleIteratorT	d_samples,
		CounterT *	d_histogram,
		int	num_levels,
		LevelT *	d_levels,
		OffsetT	num_samples,
		cudaStream_t	stream = `0`,
		bool	debug_synchronous = `false`
	)

inlinestatic

Computes an intensity histogram from a sequence of data samples using the specified bin boundary levels.

The number of histogram bins is (num_levels - 1)
The value range for bin_i is [level[i], level[i+1])
When d_temp_storage is NULL, no work is done and the required allocation size is returned in temp_storage_bytes.

Snippet: The code snippet below illustrates the computation of an six-bin histogram from a sequence of float samples

: #include <cub/cub.cuh> // or equivalently <cub/device/device_histogram.cuh>

// Declare, allocate, and initialize device-accessible pointers for input samples and

// output histogram

int num_samples; // e.g., 10

float* d_samples; // e.g., [2.2, 6.0, 7.1, 2.9, 3.5, 0.3, 2.9, 2.0, 6.1, 999.5]

int* d_histogram; // e.g., [ -, -, -, -, -, -]

int num_levels // e.g., 7 (seven level boundaries for six bins)

float* d_levels; // e.g., [0.0, 2.0, 4.0, 6.0, 8.0, 12.0, 16.0]

...

// Determine temporary device storage requirements

void* d_temp_storage = NULL;

size_t temp_storage_bytes = 0;

cub::DeviceHistogram::HistogramRange(d_temp_storage, temp_storage_bytes,

d_samples, d_histogram, num_levels, d_levels, num_samples);

// Allocate temporary storage

cudaMalloc(&d_temp_storage, temp_storage_bytes);

// Compute histograms

cub::DeviceHistogram::HistogramRange(d_temp_storage, temp_storage_bytes,

d_samples, d_histogram, num_levels, d_levels, num_samples);

// d_histogram <-- [1, 5, 0, 3, 0, 0];

Template Parameters

SampleIteratorT	[inferred] Random-access input iterator type for reading input samples. (may be a simple pointer type)
CounterT	[inferred] Integer type for histogram bin counters
LevelT	[inferred] Type for specifying boundaries (levels)
OffsetT	[inferred] Signed integer type for sequence offsets, list lengths, pointer differences, etc. (Consider using 32-bit values as offsets/lengths/etc. For example, `int` will typically yeild better performance than `size_t` in 64-bit memory mode.)

The sample value type of the input iterator

Parameters

[in]	d_temp_storage	Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to `temp_storage_bytes` and no work is done.
[in,out]	temp_storage_bytes	Reference to size in bytes of `d_temp_storage` allocation
[in]	d_samples	The pointer to the input sequence of data samples.
[out]	d_histogram	The pointer to the histogram counter output array of length `num_levels` - 1.
[in]	num_levels	The number of boundaries (levels) for delineating histogram samples. Implies that the number of bins is `num_levels` - 1.
[in]	d_levels	The pointer to the array of boundaries (levels). Bin ranges are defined by consecutive boundary pairings: lower sample value boundaries are inclusive and upper sample value boundaries are exclusive.
[in]	num_samples	The number of data samples per row in the region of interest
[in]	stream	[optional] CUDA stream to launch kernels within. Default is stream₀.
[in]	debug_synchronous	[optional] Whether or not to synchronize the stream after every kernel launch to check for errors. May cause significant slowdown. Default is `false`.

template<typename SampleIteratorT , typename CounterT , typename LevelT , typename OffsetT >

static CUB_RUNTIME_FUNCTION cudaError_t cub::DeviceHistogram::HistogramRange	(	void *	d_temp_storage,
		size_t &	temp_storage_bytes,
		SampleIteratorT	d_samples,
		CounterT *	d_histogram,
		int	num_levels,
		LevelT *	d_levels,
		OffsetT	num_row_samples,
		OffsetT	num_rows,
		size_t	row_stride_bytes,
		cudaStream_t	stream = `0`,
		bool	debug_synchronous = `false`
	)

inlinestatic

Computes an intensity histogram from a sequence of data samples using the specified bin boundary levels.

A two-dimensional region of interest within d_samples can be specified using the num_row_samples, num_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(SampleT)) == 0.
The number of histogram bins is (num_levels - 1)
The value range for bin_i is [level[i], level[i+1])
When d_temp_storage is NULL, no work is done and the required allocation size is returned in temp_storage_bytes.

Snippet: The code snippet below illustrates the computation of a six-bin histogram from a 2x5 region of interest within a flattened 2x7 array of float samples.

: #include <cub/cub.cuh> // or equivalently <cub/device/device_histogram.cuh>

// Declare, allocate, and initialize device-accessible pointers for input samples and

// output histogram

int num_row_samples; // e.g., 5

int num_rows; // e.g., 2;

int row_stride_bytes; // e.g., 7 * sizeof(float)

float* d_samples; // e.g., [2.2, 6.0, 7.1, 2.9, 3.5, -, -,

// 0.3, 2.9, 2.0, 6.1, 999.5, -, -]

int* d_histogram; // e.g., [ -, -, -, -, -, -]

int num_levels // e.g., 7 (seven level boundaries for six bins)

float *d_levels; // e.g., [0.0, 2.0, 4.0, 6.0, 8.0, 12.0, 16.0]

...

// Determine temporary device storage requirements

void* d_temp_storage = NULL;

size_t temp_storage_bytes = 0;

cub::DeviceHistogram::HistogramRange(d_temp_storage, temp_storage_bytes,

d_samples, d_histogram, num_levels, d_levels,

num_row_samples, num_rows, row_stride_bytes);

// Allocate temporary storage

cudaMalloc(&d_temp_storage, temp_storage_bytes);

// Compute histograms

cub::DeviceHistogram::HistogramRange(d_temp_storage, temp_storage_bytes,

d_samples, d_histogram, num_levels, d_levels,

num_row_samples, num_rows, row_stride_bytes);

// d_histogram <-- [1, 5, 0, 3, 0, 0];

Template Parameters

SampleIteratorT	[inferred] Random-access input iterator type for reading input samples. (may be a simple pointer type)
CounterT	[inferred] Integer type for histogram bin counters
LevelT	[inferred] Type for specifying boundaries (levels)
OffsetT	[inferred] Signed integer type for sequence offsets, list lengths, pointer differences, etc. (Consider using 32-bit values as offsets/lengths/etc. For example, `int` will typically yeild better performance than `size_t` in 64-bit memory mode.)

Parameters

[in]	d_temp_storage	Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to `temp_storage_bytes` and no work is done.
[in,out]	temp_storage_bytes	Reference to size in bytes of `d_temp_storage` allocation
[in]	d_samples	The pointer to the input sequence of data samples.
[out]	d_histogram	The pointer to the histogram counter output array of length `num_levels` - 1.
[in]	num_levels	The number of boundaries (levels) for delineating histogram samples. Implies that the number of bins is `num_levels` - 1.
[in]	d_levels	The pointer to the array of boundaries (levels). Bin ranges are defined by consecutive boundary pairings: lower sample value boundaries are inclusive and upper sample value boundaries are exclusive.
[in]	num_row_samples	The number of data samples per row in the region of interest
[in]	num_rows	The number of rows in the region of interest
[in]	row_stride_bytes	The number of bytes between starts of consecutive rows in the region of interest
[in]	stream	[optional] CUDA stream to launch kernels within. Default is stream₀.
[in]	debug_synchronous	[optional] Whether or not to synchronize the stream after every kernel launch to check for errors. May cause significant slowdown. Default is `false`.

template<int NUM_CHANNELS, int NUM_ACTIVE_CHANNELS, typename SampleIteratorT , typename CounterT , typename LevelT , typename OffsetT >

static CUB_RUNTIME_FUNCTION cudaError_t cub::DeviceHistogram::MultiHistogramRange	(	void *	d_temp_storage,
		size_t &	temp_storage_bytes,
		SampleIteratorT	d_samples,
		CounterT *	d_histogram[NUM_ACTIVE_CHANNELS],
		int	num_levels[NUM_ACTIVE_CHANNELS],
		LevelT *	d_levels[NUM_ACTIVE_CHANNELS],
		OffsetT	num_pixels,
		cudaStream_t	stream = `0`,
		bool	debug_synchronous = `false`
	)

inlinestatic

Computes per-channel intensity histograms from a sequence of multi-channel "pixel" data samples using the specified bin boundary levels.

The input is a sequence of pixel structures, where each pixel comprises a record of NUM_CHANNELS consecutive data samples (e.g., an RGBA pixel).
Of the NUM_CHANNELS specified, the function will only compute histograms for the first NUM_ACTIVE_CHANNELS (e.g., RGB histograms from RGBA pixel samples).
The number of histogram bins for channel_i is num_levels[i] - 1.
For channel_i, the range of values for all histogram bins have the same width: (upper_level[i] - lower_level[i]) / ( num_levels[i] - 1)
When d_temp_storage is NULL, no work is done and the required allocation size is returned in temp_storage_bytes.

Snippet: The code snippet below illustrates the computation of three 4-bin RGB histograms from a quad-channel sequence of RGBA pixels (8 bits per channel per pixel)

: #include <cub/cub.cuh> // or equivalently <cub/device/device_histogram.cuh>

// Declare, allocate, and initialize device-accessible pointers for input samples

// and output histograms

int num_pixels; // e.g., 5

unsigned char *d_samples; // e.g., [(2, 6, 7, 5),(3, 0, 2, 1),(7, 0, 6, 2),

// (0, 6, 7, 5),(3, 0, 2, 6)]

unsigned int *d_histogram[3]; // e.g., [[ -, -, -, -],[ -, -, -, -],[ -, -, -, -]];

int num_levels[3]; // e.g., {5, 5, 5};

unsigned int *d_levels[3]; // e.g., [ [0, 2, 4, 6, 8],

// [0, 2, 4, 6, 8],

// [0, 2, 4, 6, 8] ];

...

// Determine temporary device storage requirements

void* d_temp_storage = NULL;

size_t temp_storage_bytes = 0;

cub::DeviceHistogram::MultiHistogramRange<4, 3>(d_temp_storage, temp_storage_bytes,

d_samples, d_histogram, num_levels, d_levels, num_pixels);

// Allocate temporary storage

cudaMalloc(&d_temp_storage, temp_storage_bytes);

// Compute histograms

cub::DeviceHistogram::MultiHistogramRange<4, 3>(d_temp_storage, temp_storage_bytes,

d_samples, d_histogram, num_levels, d_levels, num_pixels);

// d_histogram <-- [ [1, 3, 0, 1],

// [3, 0, 0, 2],

// [0, 2, 0, 3] ]

Template Parameters

NUM_CHANNELS	Number of channels interleaved in the input data (may be greater than the number of channels being actively histogrammed)
NUM_ACTIVE_CHANNELS	[inferred] Number of channels actively being histogrammed
SampleIteratorT	[inferred] Random-access input iterator type for reading input samples. (may be a simple pointer type)
CounterT	[inferred] Integer type for histogram bin counters
LevelT	[inferred] Type for specifying boundaries (levels)
OffsetT	[inferred] Signed integer type for sequence offsets, list lengths, pointer differences, etc. (Consider using 32-bit values as offsets/lengths/etc. For example, `int` will typically yeild better performance than `size_t` in 64-bit memory mode.)

The sample value type of the input iterator

Parameters

[in]	d_temp_storage	Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to `temp_storage_bytes` and no work is done.
[in,out]	temp_storage_bytes	Reference to size in bytes of `d_temp_storage` allocation
[in]	d_samples	The pointer to the multi-channel input sequence of data samples. The samples from different channels are assumed to be interleaved (e.g., an array of 32-bit pixels where each pixel consists of four RGBA 8-bit samples).
[out]	d_histogram	The pointers to the histogram counter output arrays, one for each active channel. For channel_i, the allocation length of `d_histogram[i]` should be `num_levels[i]` - 1.
[in]	num_levels	The number of boundaries (levels) for delineating histogram samples in each active channel. Implies that the number of bins for channel_i is `num_levels[i]` - 1.
[in]	d_levels	The pointers to the arrays of boundaries (levels), one for each active channel. Bin ranges are defined by consecutive boundary pairings: lower sample value boundaries are inclusive and upper sample value boundaries are exclusive.
[in]	num_pixels	The number of multi-channel pixels (i.e., the length of `d_samples` / NUM_CHANNELS)
[in]	stream	[optional] CUDA stream to launch kernels within. Default is stream₀.
[in]	debug_synchronous	[optional] Whether or not to synchronize the stream after every kernel launch to check for errors. May cause significant slowdown. Default is `false`.

template<int NUM_CHANNELS, int NUM_ACTIVE_CHANNELS, typename SampleIteratorT , typename CounterT , typename LevelT , typename OffsetT >

static CUB_RUNTIME_FUNCTION cudaError_t cub::DeviceHistogram::MultiHistogramRange	(	void *	d_temp_storage,
		size_t &	temp_storage_bytes,
		SampleIteratorT	d_samples,
		CounterT *	d_histogram[NUM_ACTIVE_CHANNELS],
		int	num_levels[NUM_ACTIVE_CHANNELS],
		LevelT *	d_levels[NUM_ACTIVE_CHANNELS],
		OffsetT	num_row_pixels,
		OffsetT	num_rows,
		size_t	row_stride_bytes,
		cudaStream_t	stream = `0`,
		bool	debug_synchronous = `false`
	)

inlinestatic

Computes per-channel intensity histograms from a sequence of multi-channel "pixel" data samples using the specified bin boundary levels.

The input is a sequence of pixel structures, where each pixel comprises a record of NUM_CHANNELS consecutive data samples (e.g., an RGBA pixel).
Of the NUM_CHANNELS specified, the function will only compute histograms for the first NUM_ACTIVE_CHANNELS (e.g., RGB histograms from RGBA pixel samples).
A two-dimensional region of interest within d_samples can be specified using the num_row_samples, num_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(SampleT)) == 0.
The number of histogram bins for channel_i is num_levels[i] - 1.
For channel_i, the range of values for all histogram bins have the same width: (upper_level[i] - lower_level[i]) / ( num_levels[i] - 1)
When d_temp_storage is NULL, no work is done and the required allocation size is returned in temp_storage_bytes.

Snippet: The code snippet below illustrates the computation of three 4-bin RGB histograms from a 2x3 region of interest of within a flattened 2x4 array of quad-channel RGBA pixels (8 bits per channel per pixel).

: #include <cub/cub.cuh> // or equivalently <cub/device/device_histogram.cuh>

// Declare, allocate, and initialize device-accessible pointers for input samples

// and output histograms

int num_row_pixels; // e.g., 3

int num_rows; // e.g., 2

size_t row_stride_bytes; // e.g., 4 * sizeof(unsigned char) * NUM_CHANNELS

unsigned char* d_samples; // e.g., [(2, 6, 7, 5),(3, 0, 2, 1),(1, 1, 1, 1),(-, -, -, -),

// (7, 0, 6, 2),(0, 6, 7, 5),(3, 0, 2, 6),(-, -, -, -)]

int* d_histogram[3]; // e.g., [[ -, -, -, -],[ -, -, -, -],[ -, -, -, -]];

int num_levels[3]; // e.g., {5, 5, 5};

unsigned int* d_levels[3]; // e.g., [ [0, 2, 4, 6, 8],

// [0, 2, 4, 6, 8],

// [0, 2, 4, 6, 8] ];

...

// Determine temporary device storage requirements

void* d_temp_storage = NULL;

size_t temp_storage_bytes = 0;

cub::DeviceHistogram::MultiHistogramRange<4, 3>(d_temp_storage, temp_storage_bytes,

d_samples, d_histogram, num_levels, d_levels, num_row_pixels, num_rows, row_stride_bytes);

// Allocate temporary storage

cudaMalloc(&d_temp_storage, temp_storage_bytes);

// Compute histograms

cub::DeviceHistogram::MultiHistogramRange<4, 3>(d_temp_storage, temp_storage_bytes,

d_samples, d_histogram, num_levels, d_levels, num_row_pixels, num_rows, row_stride_bytes);

// d_histogram <-- [ [2, 3, 0, 1],

// [3, 0, 0, 2],

// [1, 2, 0, 3] ]

Template Parameters

NUM_CHANNELS	Number of channels interleaved in the input data (may be greater than the number of channels being actively histogrammed)
NUM_ACTIVE_CHANNELS	[inferred] Number of channels actively being histogrammed
SampleIteratorT	[inferred] Random-access input iterator type for reading input samples. (may be a simple pointer type)
CounterT	[inferred] Integer type for histogram bin counters
LevelT	[inferred] Type for specifying boundaries (levels)
OffsetT	[inferred] Signed integer type for sequence offsets, list lengths, pointer differences, etc. (Consider using 32-bit values as offsets/lengths/etc. For example, `int` will typically yeild better performance than `size_t` in 64-bit memory mode.)

The sample value type of the input iterator

Parameters

[in]	d_temp_storage	Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to `temp_storage_bytes` and no work is done.
[in,out]	temp_storage_bytes	Reference to size in bytes of `d_temp_storage` allocation
[in]	d_samples	The pointer to the multi-channel input sequence of data samples. The samples from different channels are assumed to be interleaved (e.g., an array of 32-bit pixels where each pixel consists of four RGBA 8-bit samples).
[out]	d_histogram	The pointers to the histogram counter output arrays, one for each active channel. For channel_i, the allocation length of `d_histogram[i]` should be `num_levels[i]` - 1.
[in]	num_levels	The number of boundaries (levels) for delineating histogram samples in each active channel. Implies that the number of bins for channel_i is `num_levels[i]` - 1.
[in]	d_levels	The pointers to the arrays of boundaries (levels), one for each active channel. Bin ranges are defined by consecutive boundary pairings: lower sample value boundaries are inclusive and upper sample value boundaries are exclusive.
[in]	num_row_pixels	The number of multi-channel pixels per row in the region of interest
[in]	num_rows	The number of rows in the region of interest
[in]	row_stride_bytes	The number of bytes between starts of consecutive rows in the region of interest
[in]	stream	[optional] CUDA stream to launch kernels within. Default is stream₀.
[in]	debug_synchronous	[optional] Whether or not to synchronize the stream after every kernel launch to check for errors. May cause significant slowdown. Default is `false`.

The documentation for this struct was generated from the following file:

device_histogram.cuh

Detailed description

Static Public Methods

Member Function Documentation