Spaces:

CVPR
/

LIVE

Runtime error

LIVE / thrust /dependencies /cub /test /test_device_histogram.cu

Xu Ma

upload all files

28958dc over 3 years ago

71 kB

	/******************************************************************************
	* Copyright (c) 2011, Duane Merrill. All rights reserved.
	* Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved.
	*
	* Redistribution and use in source and binary forms, with or without
	* modification, are permitted provided that the following conditions are met:
	* * Redistributions of source code must retain the above copyright
	* notice, this list of conditions and the following disclaimer.
	* * Redistributions in binary form must reproduce the above copyright
	* notice, this list of conditions and the following disclaimer in the
	* documentation and/or other materials provided with the distribution.
	* * Neither the name of the NVIDIA CORPORATION nor the
	* names of its contributors may be used to endorse or promote products
	* derived from this software without specific prior written permission.
	*
	* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
	* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
	* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
	* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
	* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
	* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
	* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
	* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
	* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
	* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
	*
	******************************************************************************/

	/******************************************************************************
	* Test of DeviceHistogram utilities
	******************************************************************************/

	// Ensure printing of CUDA runtime errors to console
	#define CUB_STDERR

	#include <stdio.h>
	#include <limits>
	#include <algorithm>
	#include <typeinfo>

	#if defined(QUICK_TEST) \|\| defined(QUICKER_TEST)
	#include <npp.h>
	#endif

	#include <cub/util_allocator.cuh>
	#include <cub/iterator/constant_input_iterator.cuh>
	#include <cub/device/device_histogram.cuh>

	#include "test_util.h"

	using namespace cub;


	//---------------------------------------------------------------------
	// Globals, constants and typedefs
	//---------------------------------------------------------------------


	// Dispatch types
	enum Backend
	{
	CUB, // CUB method
	NPP, // NPP method
	CDP, // GPU-based (dynamic parallelism) dispatch to CUB method
	};


	bool g_verbose_input = false;
	bool g_verbose = false;
	int g_timing_iterations = 0;
	int g_repeat = 0;
	CachingDeviceAllocator g_allocator(true);




	//---------------------------------------------------------------------
	// Dispatch to NPP histogram
	//---------------------------------------------------------------------

	#if defined(QUICK_TEST) \|\| defined(QUICKER_TEST)

	/**
	* Dispatch to single-channel 8b NPP histo-even
	*/
	template <typename CounterT, typename LevelT, typename OffsetT>
	//CUB_RUNTIME_FUNCTION __forceinline__
	cudaError_t DispatchEven(
	Int2Type<1> num_channels,
	Int2Type<1> num_active_channels,
	Int2Type<NPP> dispatch_to,
	int timing_timing_iterations,
	size_t *d_temp_storage_bytes,
	cudaError_t *d_cdp_error,

	void* d_temp_storage,
	size_t& temp_storage_bytes,
	unsigned char *d_samples, ///< [in] 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).
	CounterT *d_histogram[1], ///< [out] The pointers to the histogram counter output arrays, one for each active channel. For channel<sub><em>i</em></sub>, the allocation length of <tt>d_histograms[i]</tt> should be <tt>num_levels[i]</tt> - 1.
	int num_levels[1], ///< [in] The number of boundaries (levels) for delineating histogram samples in each active channel. Implies that the number of bins for channel<sub><em>i</em></sub> is <tt>num_levels[i]</tt> - 1.
	LevelT lower_level[1], ///< [in] The lower sample value bound (inclusive) for the lowest histogram bin in each active channel.
	LevelT upper_level[1], ///< [in] The upper sample value bound (exclusive) for the highest histogram bin in each active channel.
	OffsetT num_row_pixels, ///< [in] The number of multi-channel pixels per row in the region of interest
	OffsetT num_rows, ///< [in] The number of rows in the region of interest
	OffsetT row_stride_bytes, ///< [in] The number of bytes between starts of consecutive rows in the region of interest
	cudaStream_t stream,
	bool debug_synchronous)
	{
	typedef unsigned char SampleT;

	cudaError_t error = cudaSuccess;
	NppiSize oSizeROI = {
	num_row_pixels,
	num_rows
	};

	if (d_temp_storage_bytes == NULL)
	{
	int nDeviceBufferSize;
	nppiHistogramEvenGetBufferSize_8u_C1R(oSizeROI, num_levels[0] ,&nDeviceBufferSize);
	temp_storage_bytes = nDeviceBufferSize;
	}
	else
	{
	for (int i = 0; i < timing_timing_iterations; ++i)
	{
	// compute the histogram
	nppiHistogramEven_8u_C1R(
	d_samples,
	row_stride_bytes,
	oSizeROI,
	d_histogram[0],
	num_levels[0],
	lower_level[0],
	upper_level[0],
	(Npp8u*) d_temp_storage);
	}
	}

	return error;
	}


	/**
	* Dispatch to 3/4 8b NPP histo-even
	*/
	template <typename CounterT, typename LevelT, typename OffsetT>
	//CUB_RUNTIME_FUNCTION __forceinline__
	cudaError_t DispatchEven(
	Int2Type<4> num_channels,
	Int2Type<3> num_active_channels,
	Int2Type<NPP> dispatch_to,
	int timing_timing_iterations,
	size_t *d_temp_storage_bytes,
	cudaError_t *d_cdp_error,

	void* d_temp_storage,
	size_t& temp_storage_bytes,
	unsigned char *d_samples, ///< [in] 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).
	CounterT *d_histogram[3], ///< [out] The pointers to the histogram counter output arrays, one for each active channel. For channel<sub><em>i</em></sub>, the allocation length of <tt>d_histograms[i]</tt> should be <tt>num_levels[i]</tt> - 1.
	int num_levels[3], ///< [in] The number of boundaries (levels) for delineating histogram samples in each active channel. Implies that the number of bins for channel<sub><em>i</em></sub> is <tt>num_levels[i]</tt> - 1.
	LevelT lower_level[3], ///< [in] The lower sample value bound (inclusive) for the lowest histogram bin in each active channel.
	LevelT upper_level[3], ///< [in] The upper sample value bound (exclusive) for the highest histogram bin in each active channel.
	OffsetT num_row_pixels, ///< [in] The number of multi-channel pixels per row in the region of interest
	OffsetT num_rows, ///< [in] The number of rows in the region of interest
	OffsetT row_stride_bytes, ///< [in] The number of bytes between starts of consecutive rows in the region of interest
	cudaStream_t stream,
	bool debug_synchronous)
	{
	typedef unsigned char SampleT;

	cudaError_t error = cudaSuccess;
	NppiSize oSizeROI = {
	num_row_pixels,
	num_rows
	};

	if (d_temp_storage_bytes == NULL)
	{
	int nDeviceBufferSize;
	nppiHistogramEvenGetBufferSize_8u_AC4R(oSizeROI, num_levels ,&nDeviceBufferSize);
	temp_storage_bytes = nDeviceBufferSize;
	}
	else
	{
	for (int i = 0; i < timing_timing_iterations; ++i)
	{
	// compute the histogram
	nppiHistogramEven_8u_AC4R(
	d_samples,
	row_stride_bytes,
	oSizeROI,
	d_histogram,
	num_levels,
	lower_level,
	upper_level,
	(Npp8u*) d_temp_storage);
	}
	}

	return error;
	}


	#endif // #if defined(QUICK_TEST) \|\| defined(QUICKER_TEST)


	//---------------------------------------------------------------------
	// Dispatch to different DeviceHistogram entrypoints
	//---------------------------------------------------------------------

	template <int NUM_ACTIVE_CHANNELS, int NUM_CHANNELS, int BACKEND>
	struct Dispatch;

	template <int NUM_ACTIVE_CHANNELS, int NUM_CHANNELS>
	struct Dispatch<NUM_ACTIVE_CHANNELS, NUM_CHANNELS, CUB>
	{
	/**
	* Dispatch to CUB multi histogram-range entrypoint
	*/
	template <typename SampleIteratorT, typename CounterT, typename LevelT, typename OffsetT>
	//CUB_RUNTIME_FUNCTION __forceinline__
	static cudaError_t Range(
	int timing_timing_iterations,
	size_t /d_temp_storage_bytes*/,
	cudaError_t /d_cdp_error*/,

	void* d_temp_storage,
	size_t& temp_storage_bytes,
	SampleIteratorT d_samples, ///< [in] 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).
	CounterT *(&d_histogram)[NUM_ACTIVE_CHANNELS], ///< [out] The pointers to the histogram counter output arrays, one for each active channel. For channel<sub><em>i</em></sub>, the allocation length of <tt>d_histograms[i]</tt> should be <tt>num_levels[i]</tt> - 1.
	int *num_levels, ///< [in] The number of boundaries (levels) for delineating histogram samples in each active channel. Implies that the number of bins for channel<sub><em>i</em></sub> is <tt>num_levels[i]</tt> - 1.
	LevelT *(&d_levels)[NUM_ACTIVE_CHANNELS], ///< [in] 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.
	OffsetT num_row_pixels, ///< [in] The number of multi-channel pixels per row in the region of interest
	OffsetT num_rows, ///< [in] The number of rows in the region of interest
	OffsetT row_stride_bytes, ///< [in] The number of bytes between starts of consecutive rows in the region of interest
	cudaStream_t stream,
	bool debug_synchronous)
	{
	cudaError_t error = cudaSuccess;

	for (int i = 0; i < timing_timing_iterations; ++i)
	{
	error = DeviceHistogram::MultiHistogramRange<NUM_CHANNELS, NUM_ACTIVE_CHANNELS>(
	d_temp_storage,
	temp_storage_bytes,
	d_samples,
	d_histogram,
	num_levels,
	d_levels,
	num_row_pixels,
	num_rows,
	row_stride_bytes,
	stream,
	debug_synchronous);
	}
	return error;
	}


	/**
	* Dispatch to CUB multi histogram-even entrypoint
	*/
	template <typename SampleIteratorT, typename CounterT, typename LevelT, typename OffsetT>
	//CUB_RUNTIME_FUNCTION __forceinline__
	static cudaError_t Even(
	int timing_timing_iterations,
	size_t /d_temp_storage_bytes*/,
	cudaError_t /d_cdp_error*/,

	void* d_temp_storage,
	size_t& temp_storage_bytes,
	SampleIteratorT d_samples, ///< [in] 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).
	CounterT *(&d_histogram)[NUM_ACTIVE_CHANNELS], ///< [out] The pointers to the histogram counter output arrays, one for each active channel. For channel<sub><em>i</em></sub>, the allocation length of <tt>d_histograms[i]</tt> should be <tt>num_levels[i]</tt> - 1.
	int *num_levels, ///< [in] The number of boundaries (levels) for delineating histogram samples in each active channel. Implies that the number of bins for channel<sub><em>i</em></sub> is <tt>num_levels[i]</tt> - 1.
	LevelT *lower_level, ///< [in] The lower sample value bound (inclusive) for the lowest histogram bin in each active channel.
	LevelT *upper_level, ///< [in] The upper sample value bound (exclusive) for the highest histogram bin in each active channel.
	OffsetT num_row_pixels, ///< [in] The number of multi-channel pixels per row in the region of interest
	OffsetT num_rows, ///< [in] The number of rows in the region of interest
	OffsetT row_stride_bytes, ///< [in] The number of bytes between starts of consecutive rows in the region of interest
	cudaStream_t stream,
	bool debug_synchronous)
	{
	typedef typename std::iterator_traits<SampleIteratorT>::value_type SampleT;

	cudaError_t error = cudaSuccess;
	for (int i = 0; i < timing_timing_iterations; ++i)
	{
	error = DeviceHistogram::MultiHistogramEven<NUM_CHANNELS, NUM_ACTIVE_CHANNELS>(
	d_temp_storage,
	temp_storage_bytes,
	d_samples,
	d_histogram,
	num_levels,
	lower_level,
	upper_level,
	num_row_pixels,
	num_rows,
	row_stride_bytes,
	stream,
	debug_synchronous);
	}
	return error;
	}

	};


	template <>
	struct Dispatch<1, 1, CUB>
	{

	/**
	* Dispatch to CUB single histogram-range entrypoint
	*/
	template <typename SampleIteratorT, typename CounterT, typename LevelT, typename OffsetT>
	//CUB_RUNTIME_FUNCTION __forceinline__
	static cudaError_t Range(
	int timing_timing_iterations,
	size_t /d_temp_storage_bytes*/,
	cudaError_t /d_cdp_error*/,

	void* d_temp_storage,
	size_t& temp_storage_bytes,
	SampleIteratorT d_samples, ///< [in] 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).
	CounterT* (&d_histogram)[1], ///< [out] The pointers to the histogram counter output arrays, one for each active channel. For channel<sub><em>i</em></sub>, the allocation length of <tt>d_histograms[i]</tt> should be <tt>num_levels[i]</tt> - 1.
	int *num_levels, ///< [in] The number of boundaries (levels) for delineating histogram samples in each active channel. Implies that the number of bins for channel<sub><em>i</em></sub> is <tt>num_levels[i]</tt> - 1.
	LevelT (&d_levels)[1], ///< [in] 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.
	OffsetT num_row_pixels, ///< [in] The number of multi-channel pixels per row in the region of interest
	OffsetT num_rows, ///< [in] The number of rows in the region of interest
	OffsetT row_stride_bytes, ///< [in] The number of bytes between starts of consecutive rows in the region of interest
	cudaStream_t stream,
	bool debug_synchronous)
	{
	cudaError_t error = cudaSuccess;
	for (int i = 0; i < timing_timing_iterations; ++i)
	{
	error = DeviceHistogram::HistogramRange(
	d_temp_storage,
	temp_storage_bytes,
	d_samples,
	d_histogram[0],
	num_levels[0],
	d_levels[0],
	num_row_pixels,
	num_rows,
	row_stride_bytes,
	stream,
	debug_synchronous);
	}
	return error;
	}


	/**
	* Dispatch to CUB single histogram-even entrypoint
	*/
	template <typename SampleIteratorT, typename CounterT, typename LevelT, typename OffsetT>
	//CUB_RUNTIME_FUNCTION __forceinline__
	static cudaError_t Even(
	int timing_timing_iterations,
	size_t /d_temp_storage_bytes*/,
	cudaError_t /d_cdp_error*/,

	void* d_temp_storage,
	size_t& temp_storage_bytes,
	SampleIteratorT d_samples, ///< [in] 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).
	CounterT* (&d_histogram)[1], ///< [out] The pointers to the histogram counter output arrays, one for each active channel. For channel<sub><em>i</em></sub>, the allocation length of <tt>d_histograms[i]</tt> should be <tt>num_levels[i]</tt> - 1.
	int *num_levels, ///< [in] The number of boundaries (levels) for delineating histogram samples in each active channel. Implies that the number of bins for channel<sub><em>i</em></sub> is <tt>num_levels[i]</tt> - 1.
	LevelT *lower_level, ///< [in] The lower sample value bound (inclusive) for the lowest histogram bin in each active channel.
	LevelT *upper_level, ///< [in] The upper sample value bound (exclusive) for the highest histogram bin in each active channel.
	OffsetT num_row_pixels, ///< [in] The number of multi-channel pixels per row in the region of interest
	OffsetT num_rows, ///< [in] The number of rows in the region of interest
	OffsetT row_stride_bytes, ///< [in] The number of bytes between starts of consecutive rows in the region of interest
	cudaStream_t stream,
	bool debug_synchronous)
	{
	cudaError_t error = cudaSuccess;
	for (int i = 0; i < timing_timing_iterations; ++i)
	{
	error = DeviceHistogram::HistogramEven(
	d_temp_storage,
	temp_storage_bytes,
	d_samples,
	d_histogram[0],
	num_levels[0],
	lower_level[0],
	upper_level[0],
	num_row_pixels,
	num_rows,
	row_stride_bytes,
	stream,
	debug_synchronous);
	}
	return error;
	}

	};



	//---------------------------------------------------------------------
	// CUDA nested-parallelism test kernel
	//---------------------------------------------------------------------

	/**
	* Simple wrapper kernel to invoke DeviceHistogram
	* /
	template <int BINS, int NUM_CHANNELS, int NUM_ACTIVE_CHANNELS, typename SampleT, typename SampleIteratorT, typename CounterT, int ALGORITHM>
	__global__ void CnpDispatchKernel(
	Int2Type<ALGORITHM> algorithm,
	int timing_timing_iterations,
	size_t *d_temp_storage_bytes,
	cudaError_t *d_cdp_error,

	void* d_temp_storage,
	size_t temp_storage_bytes,
	SampleT *d_samples,
	SampleIteratorT d_sample_itr,
	ArrayWrapper<CounterT*, NUM_ACTIVE_CHANNELS> d_out_histograms,
	int num_samples,
	bool debug_synchronous)
	{
	#ifndef CUB_CDP
	*d_cdp_error = cudaErrorNotSupported;
	#else
	*d_cdp_error = Dispatch<BINS, NUM_CHANNELS, NUM_ACTIVE_CHANNELS>(algorithm, Int2Type<false>(), timing_timing_iterations, d_temp_storage_bytes, d_cdp_error, d_temp_storage, temp_storage_bytes, d_samples, d_sample_itr, d_out_histograms.array, num_samples, 0, debug_synchronous);
	*d_temp_storage_bytes = temp_storage_bytes;
	#endif
	}


	/ **
	* Dispatch to CDP kernel
	* /
	template <int BINS, int NUM_CHANNELS, int NUM_ACTIVE_CHANNELS, typename SampleT, typename SampleIteratorT, typename CounterT, int ALGORITHM>
	cudaError_t Dispatch(
	Int2Type<ALGORITHM> algorithm,
	Int2Type<true> use_cdp,
	int timing_timing_iterations,
	size_t *d_temp_storage_bytes,
	cudaError_t *d_cdp_error,

	void* d_temp_storage,
	size_t& temp_storage_bytes,
	SampleT *d_samples,
	SampleIteratorT d_sample_itr,
	CounterT *d_histograms[NUM_ACTIVE_CHANNELS],
	int num_samples,
	cudaStream_t stream,
	bool debug_synchronous)
	{
	// Setup array wrapper for histogram channel output (because we can't pass static arrays as kernel parameters)
	ArrayWrapper<CounterT*, NUM_ACTIVE_CHANNELS> d_histo_wrapper;
	for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL)
	d_histo_wrapper.array[CHANNEL] = d_histograms[CHANNEL];

	// Invoke kernel to invoke device-side dispatch
	CnpDispatchKernel<BINS, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleIteratorT, CounterT, ALGORITHM><<<1,1>>>(algorithm, timing_timing_iterations, d_temp_storage_bytes, d_cdp_error, d_temp_storage, temp_storage_bytes, d_samples, d_sample_itr, d_histo_wrapper, num_samples, debug_synchronous);

	// Copy out temp_storage_bytes
	CubDebugExit(cudaMemcpy(&temp_storage_bytes, d_temp_storage_bytes, sizeof(size_t) * 1, cudaMemcpyDeviceToHost));

	// Copy out error
	cudaError_t retval;
	CubDebugExit(cudaMemcpy(&retval, d_cdp_error, sizeof(cudaError_t) * 1, cudaMemcpyDeviceToHost));
	return retval;
	}
	*/


	//---------------------------------------------------------------------
	// Test generation
	//---------------------------------------------------------------------

	// Searches for bin given a list of bin-boundary levels
	template <typename LevelT>
	struct SearchTransform
	{
	LevelT *levels; // Pointer to levels array
	int num_levels; // Number of levels in array

	// Functor for converting samples to bin-ids (num_levels is returned if sample is out of range)
	template <typename SampleT>
	int operator()(SampleT sample)
	{
	int bin = int(std::upper_bound(levels, levels + num_levels, (LevelT) sample) - levels - 1);
	if (bin < 0)
	{
	// Sample out of range
	return num_levels;
	}
	return bin;
	}
	};


	// Scales samples to evenly-spaced bins
	template <typename LevelT>
	struct ScaleTransform
	{
	int num_levels; // Number of levels in array
	LevelT max; // Max sample level (exclusive)
	LevelT min; // Min sample level (inclusive)
	LevelT scale; // Bin scaling factor

	void Init(
	int num_levels, // Number of levels in array
	LevelT max, // Max sample level (exclusive)
	LevelT min, // Min sample level (inclusive)
	LevelT scale) // Bin scaling factor
	{
	this->num_levels = num_levels;
	this->max = max;
	this->min = min;
	this->scale = scale;
	}

	// Functor for converting samples to bin-ids (num_levels is returned if sample is out of range)
	template <typename SampleT>
	int operator()(SampleT sample)
	{
	if ((sample < min) \|\| (sample >= max))
	{
	// Sample out of range
	return num_levels;
	}

	return (int) ((((LevelT) sample) - min) / scale);
	}
	};

	// Scales samples to evenly-spaced bins
	template <>
	struct ScaleTransform<float>
	{
	int num_levels; // Number of levels in array
	float max; // Max sample level (exclusive)
	float min; // Min sample level (inclusive)
	float scale; // Bin scaling factor

	void Init(
	int num_levels, // Number of levels in array
	float max, // Max sample level (exclusive)
	float min, // Min sample level (inclusive)
	float scale) // Bin scaling factor
	{
	this->num_levels = num_levels;
	this->max = max;
	this->min = min;
	this->scale = 1.0f / scale;
	}

	// Functor for converting samples to bin-ids (num_levels is returned if sample is out of range)
	template <typename SampleT>
	int operator()(SampleT sample)
	{
	if ((sample < min) \|\| (sample >= max))
	{
	// Sample out of range
	return num_levels;
	}

	return (int) ((((float) sample) - min) * scale);
	}
	};


	/**
	* Generate sample
	*/
	template <typename T, typename LevelT>
	void Sample(T &datum, LevelT max_level, int entropy_reduction)
	{
	unsigned int max = (unsigned int) -1;
	unsigned int bits;
	RandomBits(bits, entropy_reduction);
	float fraction = (float(bits) / max);

	datum = (T) (fraction * max_level);
	}


	/**
	* Initialize histogram samples
	*/
	template <
	int NUM_CHANNELS,
	int NUM_ACTIVE_CHANNELS,
	typename LevelT,
	typename SampleT,
	typename OffsetT>
	void InitializeSamples(
	LevelT max_level,
	int entropy_reduction,
	SampleT *h_samples,
	OffsetT num_row_pixels, ///< [in] The number of multi-channel pixels per row in the region of interest
	OffsetT num_rows, ///< [in] The number of rows in the region of interest
	OffsetT row_stride_bytes) ///< [in] The number of bytes between starts of consecutive rows in the region of interest
	{
	// Initialize samples
	for (OffsetT row = 0; row < num_rows; ++row)
	{
	for (OffsetT pixel = 0; pixel < num_row_pixels; ++pixel)
	{
	for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel)
	{
	// Sample offset
	OffsetT offset = (row * (row_stride_bytes / sizeof(SampleT))) + (pixel * NUM_CHANNELS) + channel;

	// Init sample value
	Sample(h_samples[offset], max_level, entropy_reduction);
	if (g_verbose_input)
	{
	if (channel > 0) printf(", ");
	std::cout << CoutCast(h_samples[offset]);
	}
	}
	}
	}
	}


	/**
	* Initialize histogram solutions
	*/
	template <
	int NUM_CHANNELS,
	int NUM_ACTIVE_CHANNELS,
	typename CounterT,
	typename SampleIteratorT,
	typename TransformOp,
	typename OffsetT>
	void InitializeBins(
	SampleIteratorT h_samples,
	int num_levels[NUM_ACTIVE_CHANNELS], ///< [in] The number of boundaries (levels) for delineating histogram samples in each active channel. Implies that the number of bins for channel<sub><em>i</em></sub> is <tt>num_levels[i]</tt> - 1.
	TransformOp transform_op[NUM_ACTIVE_CHANNELS], ///< [in] The lower sample value bound (inclusive) for the lowest histogram bin in each active channel.
	CounterT *h_histogram[NUM_ACTIVE_CHANNELS], ///< [out] The pointers to the histogram counter output arrays, one for each active channel. For channel<sub><em>i</em></sub>, the allocation length of <tt>d_histograms[i]</tt> should be <tt>num_levels[i]</tt> - 1.
	OffsetT num_row_pixels, ///< [in] The number of multi-channel pixels per row in the region of interest
	OffsetT num_rows, ///< [in] The number of rows in the region of interest
	OffsetT row_stride_bytes) ///< [in] The number of bytes between starts of consecutive rows in the region of interest
	{
	typedef typename std::iterator_traits<SampleIteratorT>::value_type SampleT;

	// Init bins
	for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL)
	{
	for (int bin = 0; bin < num_levels[CHANNEL] - 1; ++bin)
	{
	h_histogram[CHANNEL][bin] = 0;
	}
	}

	// Initialize samples
	if (g_verbose_input) printf("Samples: \n");
	for (OffsetT row = 0; row < num_rows; ++row)
	{
	for (OffsetT pixel = 0; pixel < num_row_pixels; ++pixel)
	{
	if (g_verbose_input) printf("[");
	for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel)
	{
	// Sample offset
	OffsetT offset = (row * (row_stride_bytes / sizeof(SampleT))) + (pixel * NUM_CHANNELS) + channel;

	// Update sample bin
	int bin = transform_op[channel](h_samples[offset]);
	if (g_verbose_input) printf(" (%d)", bin); fflush(stdout);
	if ((bin >= 0) && (bin < num_levels[channel] - 1))
	{
	// valid bin
	h_histogram[channel][bin]++;
	}
	}
	if (g_verbose_input) printf("]");
	}
	if (g_verbose_input) printf("\n\n");
	}
	}



	/**
	* Test histogram-even
	*/
	template <
	Backend BACKEND,
	int NUM_CHANNELS,
	int NUM_ACTIVE_CHANNELS,
	typename SampleT,
	typename CounterT,
	typename LevelT,
	typename OffsetT,
	typename SampleIteratorT>
	void TestEven(
	LevelT max_level,
	int entropy_reduction,
	int num_levels[NUM_ACTIVE_CHANNELS], ///< [in] The number of boundaries (levels) for delineating histogram samples in each active channel. Implies that the number of bins for channel<sub><em>i</em></sub> is <tt>num_levels[i]</tt> - 1.
	LevelT lower_level[NUM_ACTIVE_CHANNELS], ///< [in] The lower sample value bound (inclusive) for the lowest histogram bin in each active channel.
	LevelT upper_level[NUM_ACTIVE_CHANNELS], ///< [in] The upper sample value bound (exclusive) for the highest histogram bin in each active channel.
	OffsetT num_row_pixels, ///< [in] The number of multi-channel pixels per row in the region of interest
	OffsetT num_rows, ///< [in] The number of rows in the region of interest
	OffsetT row_stride_bytes, ///< [in] The number of bytes between starts of consecutive rows in the region of interest
	SampleIteratorT h_samples,
	SampleIteratorT d_samples)
	{
	OffsetT total_samples = num_rows * (row_stride_bytes / sizeof(SampleT));

	printf("\n----------------------------\n");
	printf("%s cub::DeviceHistogramEven (%s) %d pixels (%d height, %d width, %d-byte row stride), %d %d-byte %s samples (entropy reduction %d), %s counters, %d/%d channels, max sample ",
	(BACKEND == CDP) ? "CDP CUB" : (BACKEND == NPP) ? "NPP" : "CUB",
	(IsPointer<SampleIteratorT>::VALUE) ? "pointer" : "iterator",
	(int) (num_row_pixels * num_rows),
	(int) num_rows,
	(int) num_row_pixels,
	(int) row_stride_bytes,
	(int) total_samples,
	(int) sizeof(SampleT),
	typeid(SampleT).name(),
	entropy_reduction,
	typeid(CounterT).name(),
	NUM_ACTIVE_CHANNELS,
	NUM_CHANNELS);
	std::cout << CoutCast(max_level) << "\n";
	for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel)
	std::cout << "\n\tChannel " << channel << ": " << num_levels[channel] - 1 << " bins [" << lower_level[channel] << ", " << upper_level[channel] << ")\n";
	fflush(stdout);

	// Allocate and initialize host and device data

	typedef SampleT Foo; // rename type to quelch gcc warnings (bug?)
	CounterT* h_histogram[NUM_ACTIVE_CHANNELS];
	ScaleTransform<LevelT> transform_op[NUM_ACTIVE_CHANNELS];

	for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel)
	{
	int bins = num_levels[channel] - 1;
	h_histogram[channel] = new CounterT[bins];

	transform_op[channel].Init(
	num_levels[channel],
	upper_level[channel],
	lower_level[channel],
	((upper_level[channel] - lower_level[channel]) / bins));
	}

	InitializeBins<NUM_CHANNELS, NUM_ACTIVE_CHANNELS>(
	h_samples, num_levels, transform_op, h_histogram, num_row_pixels, num_rows, row_stride_bytes);

	// Allocate and initialize device data

	CounterT* d_histogram[NUM_ACTIVE_CHANNELS];
	for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel)
	{
	CubDebugExit(g_allocator.DeviceAllocate((void*)&d_histogram[channel], sizeof(CounterT) (num_levels[channel] - 1)));
	CubDebugExit(cudaMemset(d_histogram[channel], 0, sizeof(CounterT) * (num_levels[channel] - 1)));
	}

	// Allocate CDP device arrays
	size_t *d_temp_storage_bytes = NULL;
	cudaError_t *d_cdp_error = NULL;
	CubDebugExit(g_allocator.DeviceAllocate((void*)&d_temp_storage_bytes, sizeof(size_t) 1));
	CubDebugExit(g_allocator.DeviceAllocate((void*)&d_cdp_error, sizeof(cudaError_t) 1));

	// Allocate temporary storage
	void *d_temp_storage = NULL;
	size_t temp_storage_bytes = 0;

	Dispatch<NUM_ACTIVE_CHANNELS, NUM_CHANNELS, BACKEND>::Even(
	1, d_temp_storage_bytes, d_cdp_error,
	d_temp_storage, temp_storage_bytes,
	d_samples, d_histogram, num_levels, lower_level, upper_level,
	num_row_pixels, num_rows, row_stride_bytes,
	0, true);

	// Allocate temporary storage with "canary" zones
	int canary_bytes = 256;
	char canary_token = 8;
	char* canary_zone = new char[canary_bytes];

	memset(canary_zone, canary_token, canary_bytes);
	CubDebugExit(g_allocator.DeviceAllocate(&d_temp_storage, temp_storage_bytes + (canary_bytes * 2)));
	CubDebugExit(cudaMemset(d_temp_storage, canary_token, temp_storage_bytes + (canary_bytes * 2)));

	// Run warmup/correctness iteration
	Dispatch<NUM_ACTIVE_CHANNELS, NUM_CHANNELS, BACKEND>::Even(
	1, d_temp_storage_bytes, d_cdp_error,
	((char *) d_temp_storage) + canary_bytes, temp_storage_bytes,
	d_samples, d_histogram, num_levels, lower_level, upper_level,
	num_row_pixels, num_rows, row_stride_bytes,
	0, true);

	// Check canary zones
	int error = CompareDeviceResults(canary_zone, (char *) d_temp_storage, canary_bytes, true, g_verbose);
	AssertEquals(0, error);
	error = CompareDeviceResults(canary_zone, ((char *) d_temp_storage) + canary_bytes + temp_storage_bytes, canary_bytes, true, g_verbose);
	AssertEquals(0, error);

	// Flush any stdout/stderr
	CubDebugExit(cudaPeekAtLastError());
	CubDebugExit(cudaDeviceSynchronize());
	fflush(stdout);
	fflush(stderr);

	// Check for correctness (and display results, if specified)
	for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel)
	{
	int channel_error = CompareDeviceResults(h_histogram[channel], d_histogram[channel], num_levels[channel] - 1, true, g_verbose);
	printf("\tChannel %d %s", channel, channel_error ? "FAIL" : "PASS\n");
	error \|= channel_error;
	}

	// Performance
	GpuTimer gpu_timer;
	gpu_timer.Start();

	Dispatch<NUM_ACTIVE_CHANNELS, NUM_CHANNELS, BACKEND>::Even(
	g_timing_iterations, d_temp_storage_bytes, d_cdp_error,
	((char *) d_temp_storage) + canary_bytes, temp_storage_bytes,
	d_samples, d_histogram, num_levels, lower_level, upper_level,
	num_row_pixels, num_rows, row_stride_bytes,
	0, false);

	gpu_timer.Stop();
	float elapsed_millis = gpu_timer.ElapsedMillis();

	// Display performance
	if (g_timing_iterations > 0)
	{
	float avg_millis = elapsed_millis / g_timing_iterations;
	float giga_rate = float(total_samples) / avg_millis / 1000.0f / 1000.0f;
	float giga_bandwidth = giga_rate * sizeof(SampleT);
	printf("\t%.3f avg ms, %.3f billion samples/s, %.3f billion bins/s, %.3f billion pixels/s, %.3f logical GB/s",
	avg_millis,
	giga_rate,
	giga_rate * NUM_ACTIVE_CHANNELS / NUM_CHANNELS,
	giga_rate / NUM_CHANNELS,
	giga_bandwidth);
	}

	printf("\n\n");

	for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel)
	{
	if (h_histogram[channel])
	delete[] h_histogram[channel];

	if (d_histogram[channel])
	CubDebugExit(g_allocator.DeviceFree(d_histogram[channel]));
	}

	if (d_temp_storage_bytes) CubDebugExit(g_allocator.DeviceFree(d_temp_storage_bytes));
	if (d_cdp_error) CubDebugExit(g_allocator.DeviceFree(d_cdp_error));
	if (d_temp_storage) CubDebugExit(g_allocator.DeviceFree(d_temp_storage));

	// Correctness asserts
	AssertEquals(0, error);
	}


	/**
	* Test histogram-even (native pointer input)
	*/
	template <
	Backend BACKEND,
	int NUM_CHANNELS,
	int NUM_ACTIVE_CHANNELS,
	typename SampleT,
	typename CounterT,
	typename LevelT,
	typename OffsetT>
	void TestEvenNative(
	LevelT max_level,
	int entropy_reduction,
	int num_levels[NUM_ACTIVE_CHANNELS], ///< [in] The number of boundaries (levels) for delineating histogram samples in each active channel. Implies that the number of bins for channel<sub><em>i</em></sub> is <tt>num_levels[i]</tt> - 1.
	LevelT lower_level[NUM_ACTIVE_CHANNELS], ///< [in] The lower sample value bound (inclusive) for the lowest histogram bin in each active channel.
	LevelT upper_level[NUM_ACTIVE_CHANNELS], ///< [in] The upper sample value bound (exclusive) for the highest histogram bin in each active channel.
	OffsetT num_row_pixels, ///< [in] The number of multi-channel pixels per row in the region of interest
	OffsetT num_rows, ///< [in] The number of rows in the region of interest
	OffsetT row_stride_bytes) ///< [in] The number of bytes between starts of consecutive rows in the region of interest
	{
	OffsetT total_samples = num_rows * (row_stride_bytes / sizeof(SampleT));

	// Allocate and initialize host sample data
	typedef SampleT Foo; // rename type to quelch gcc warnings (bug?)
	SampleT* h_samples = new Foo[total_samples];

	InitializeSamples<NUM_CHANNELS, NUM_ACTIVE_CHANNELS>(
	max_level, entropy_reduction, h_samples, num_row_pixels, num_rows, row_stride_bytes);

	// Allocate and initialize device data
	SampleT* d_samples = NULL;
	CubDebugExit(g_allocator.DeviceAllocate((void*)&d_samples, sizeof(SampleT) total_samples));
	CubDebugExit(cudaMemcpy(d_samples, h_samples, sizeof(SampleT) * total_samples, cudaMemcpyHostToDevice));

	TestEven<BACKEND, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleT, CounterT, LevelT, OffsetT>(
	max_level, entropy_reduction, num_levels, lower_level, upper_level,
	num_row_pixels, num_rows, row_stride_bytes,
	h_samples, d_samples);

	// Cleanup
	if (h_samples) delete[] h_samples;
	if (d_samples) CubDebugExit(g_allocator.DeviceFree(d_samples));
	}


	/**
	* Test histogram-even (native pointer input)
	*/
	template <
	Backend BACKEND,
	int NUM_CHANNELS,
	int NUM_ACTIVE_CHANNELS,
	typename SampleT,
	typename CounterT,
	typename LevelT,
	typename OffsetT>
	void TestEvenIterator(
	LevelT max_level,
	int entropy_reduction,
	int num_levels[NUM_ACTIVE_CHANNELS], ///< [in] The number of boundaries (levels) for delineating histogram samples in each active channel. Implies that the number of bins for channel<sub><em>i</em></sub> is <tt>num_levels[i]</tt> - 1.
	LevelT lower_level[NUM_ACTIVE_CHANNELS], ///< [in] The lower sample value bound (inclusive) for the lowest histogram bin in each active channel.
	LevelT upper_level[NUM_ACTIVE_CHANNELS], ///< [in] The upper sample value bound (exclusive) for the highest histogram bin in each active channel.
	OffsetT num_row_pixels, ///< [in] The number of multi-channel pixels per row in the region of interest
	OffsetT num_rows, ///< [in] The number of rows in the region of interest
	OffsetT row_stride_bytes) ///< [in] The number of bytes between starts of consecutive rows in the region of interest
	{
	SampleT sample = (SampleT) lower_level[0];
	ConstantInputIterator<SampleT> sample_itr(sample);

	TestEven<BACKEND, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleT, CounterT, LevelT, OffsetT>(
	max_level, entropy_reduction, num_levels, lower_level, upper_level,
	num_row_pixels, num_rows, row_stride_bytes,
	sample_itr, sample_itr);

	}


	/**
	* Test histogram-range
	*/
	template <
	Backend BACKEND,
	int NUM_CHANNELS,
	int NUM_ACTIVE_CHANNELS,
	typename SampleT,
	typename CounterT,
	typename LevelT,
	typename OffsetT>
	void TestRange(
	LevelT max_level,
	int entropy_reduction,
	int num_levels[NUM_ACTIVE_CHANNELS], ///< [in] The number of boundaries (levels) for delineating histogram samples in each active channel. Implies that the number of bins for channel<sub><em>i</em></sub> is <tt>num_levels[i]</tt> - 1.
	LevelT* levels[NUM_ACTIVE_CHANNELS], ///< [in] The lower sample value bound (inclusive) for the lowest histogram bin in each active channel.
	OffsetT num_row_pixels, ///< [in] The number of multi-channel pixels per row in the region of interest
	OffsetT num_rows, ///< [in] The number of rows in the region of interest
	OffsetT row_stride_bytes) ///< [in] The number of bytes between starts of consecutive rows in the region of interest
	{
	OffsetT total_samples = num_rows * (row_stride_bytes / sizeof(SampleT));

	printf("\n----------------------------\n");
	printf("%s cub::DeviceHistogramRange %d pixels (%d height, %d width, %d-byte row stride), %d %d-byte %s samples (entropy reduction %d), %s counters, %d/%d channels, max sample ",
	(BACKEND == CDP) ? "CDP CUB" : (BACKEND == NPP) ? "NPP" : "CUB",
	(int) (num_row_pixels * num_rows),
	(int) num_rows,
	(int) num_row_pixels,
	(int) row_stride_bytes,
	(int) total_samples,
	(int) sizeof(SampleT),
	typeid(SampleT).name(),
	entropy_reduction,
	typeid(CounterT).name(),
	NUM_ACTIVE_CHANNELS,
	NUM_CHANNELS);
	std::cout << CoutCast(max_level) << "\n";
	for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel)
	{
	printf("Channel %d: %d bins [", channel, num_levels[channel] - 1);
	std::cout << levels[channel][0];
	for (int level = 1; level < num_levels[channel]; ++level)
	std::cout << ", " << levels[channel][level];
	printf("]\n");
	}
	fflush(stdout);

	// Allocate and initialize host and device data
	typedef SampleT Foo; // rename type to quelch gcc warnings (bug?)
	SampleT* h_samples = new Foo[total_samples];
	CounterT* h_histogram[NUM_ACTIVE_CHANNELS];
	SearchTransform<LevelT> transform_op[NUM_ACTIVE_CHANNELS];

	for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel)
	{
	transform_op[channel].levels = levels[channel];
	transform_op[channel].num_levels = num_levels[channel];

	int bins = num_levels[channel] - 1;
	h_histogram[channel] = new CounterT[bins];
	}

	InitializeSamples<NUM_CHANNELS, NUM_ACTIVE_CHANNELS>(
	max_level, entropy_reduction, h_samples, num_row_pixels, num_rows, row_stride_bytes);

	InitializeBins<NUM_CHANNELS, NUM_ACTIVE_CHANNELS>(
	h_samples, num_levels, transform_op, h_histogram, num_row_pixels, num_rows, row_stride_bytes);

	// Allocate and initialize device data
	SampleT* d_samples = NULL;
	LevelT* d_levels[NUM_ACTIVE_CHANNELS];
	CounterT* d_histogram[NUM_ACTIVE_CHANNELS];

	CubDebugExit(g_allocator.DeviceAllocate((void*)&d_samples, sizeof(SampleT) total_samples));
	CubDebugExit(cudaMemcpy(d_samples, h_samples, sizeof(SampleT) * total_samples, cudaMemcpyHostToDevice));

	for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel)
	{
	CubDebugExit(g_allocator.DeviceAllocate((void*)&d_levels[channel], sizeof(LevelT) num_levels[channel]));
	CubDebugExit(cudaMemcpy(d_levels[channel], levels[channel], sizeof(LevelT) * num_levels[channel], cudaMemcpyHostToDevice));

	int bins = num_levels[channel] - 1;
	CubDebugExit(g_allocator.DeviceAllocate((void*)&d_histogram[channel], sizeof(CounterT) bins));
	CubDebugExit(cudaMemset(d_histogram[channel], 0, sizeof(CounterT) * bins));
	}

	// Allocate CDP device arrays
	size_t *d_temp_storage_bytes = NULL;
	cudaError_t *d_cdp_error = NULL;

	CubDebugExit(g_allocator.DeviceAllocate((void*)&d_temp_storage_bytes, sizeof(size_t) 1));
	CubDebugExit(g_allocator.DeviceAllocate((void*)&d_cdp_error, sizeof(cudaError_t) 1));

	// Allocate temporary storage
	void *d_temp_storage = NULL;
	size_t temp_storage_bytes = 0;

	Dispatch<NUM_ACTIVE_CHANNELS, NUM_CHANNELS, BACKEND>::Range(
	1, d_temp_storage_bytes, d_cdp_error,
	d_temp_storage, temp_storage_bytes,
	d_samples,
	d_histogram,
	num_levels, d_levels,
	num_row_pixels, num_rows, row_stride_bytes,
	0, true);

	// Allocate temporary storage with "canary" zones
	int canary_bytes = 256;
	char canary_token = 9;
	char* canary_zone = new char[canary_bytes];

	memset(canary_zone, canary_token, canary_bytes);
	CubDebugExit(g_allocator.DeviceAllocate(&d_temp_storage, temp_storage_bytes + (canary_bytes * 2)));
	CubDebugExit(cudaMemset(d_temp_storage, canary_token, temp_storage_bytes + (canary_bytes * 2)));

	// Run warmup/correctness iteration
	Dispatch<NUM_ACTIVE_CHANNELS, NUM_CHANNELS, BACKEND>::Range(
	1, d_temp_storage_bytes, d_cdp_error,
	((char *) d_temp_storage) + canary_bytes, temp_storage_bytes,
	d_samples,
	d_histogram,
	num_levels, d_levels,
	num_row_pixels, num_rows, row_stride_bytes,
	0, true);

	// Check canary zones
	int error = CompareDeviceResults(canary_zone, (char *) d_temp_storage, canary_bytes, true, g_verbose);
	AssertEquals(0, error);
	error = CompareDeviceResults(canary_zone, ((char *) d_temp_storage) + canary_bytes + temp_storage_bytes, canary_bytes, true, g_verbose);
	AssertEquals(0, error);

	// Flush any stdout/stderr
	CubDebugExit(cudaPeekAtLastError());
	CubDebugExit(cudaDeviceSynchronize());
	fflush(stdout);
	fflush(stderr);

	// Check for correctness (and display results, if specified)
	for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel)
	{
	int channel_error = CompareDeviceResults(h_histogram[channel], d_histogram[channel], num_levels[channel] - 1, true, g_verbose);
	printf("\tChannel %d %s", channel, channel_error ? "FAIL" : "PASS\n");
	error \|= channel_error;
	}

	// Performance
	GpuTimer gpu_timer;
	gpu_timer.Start();

	Dispatch<NUM_ACTIVE_CHANNELS, NUM_CHANNELS, BACKEND>::Range(
	g_timing_iterations, d_temp_storage_bytes, d_cdp_error,
	((char *) d_temp_storage) + canary_bytes, temp_storage_bytes,
	d_samples,
	d_histogram,
	num_levels, d_levels,
	num_row_pixels, num_rows, row_stride_bytes,
	0, false);

	gpu_timer.Stop();
	float elapsed_millis = gpu_timer.ElapsedMillis();

	// Display performance
	if (g_timing_iterations > 0)
	{
	float avg_millis = elapsed_millis / g_timing_iterations;
	float giga_rate = float(total_samples) / avg_millis / 1000.0f / 1000.0f;
	float giga_bandwidth = giga_rate * sizeof(SampleT);
	printf("\t%.3f avg ms, %.3f billion samples/s, %.3f billion bins/s, %.3f billion pixels/s, %.3f logical GB/s",
	avg_millis,
	giga_rate,
	giga_rate * NUM_ACTIVE_CHANNELS / NUM_CHANNELS,
	giga_rate / NUM_CHANNELS,
	giga_bandwidth);
	}

	printf("\n\n");

	// Cleanup
	if (h_samples) delete[] h_samples;

	for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel)
	{
	if (h_histogram[channel])
	delete[] h_histogram[channel];

	if (d_histogram[channel])
	CubDebugExit(g_allocator.DeviceFree(d_histogram[channel]));

	if (d_levels[channel])
	CubDebugExit(g_allocator.DeviceFree(d_levels[channel]));
	}

	if (d_samples) CubDebugExit(g_allocator.DeviceFree(d_samples));
	if (d_temp_storage_bytes) CubDebugExit(g_allocator.DeviceFree(d_temp_storage_bytes));
	if (d_cdp_error) CubDebugExit(g_allocator.DeviceFree(d_cdp_error));
	if (d_temp_storage) CubDebugExit(g_allocator.DeviceFree(d_temp_storage));

	// Correctness asserts
	AssertEquals(0, error);
	}


	/**
	* Test histogram-even
	*/
	template <
	Backend BACKEND,
	typename SampleT,
	int NUM_CHANNELS,
	int NUM_ACTIVE_CHANNELS,
	typename CounterT,
	typename LevelT,
	typename OffsetT>
	void TestEven(
	OffsetT num_row_pixels,
	OffsetT num_rows,
	OffsetT row_stride_bytes,
	int entropy_reduction,
	int num_levels[NUM_ACTIVE_CHANNELS],
	LevelT max_level,
	int max_num_levels)
	{
	LevelT lower_level[NUM_ACTIVE_CHANNELS];
	LevelT upper_level[NUM_ACTIVE_CHANNELS];

	// Find smallest level increment
	int max_bins = max_num_levels - 1;
	LevelT min_level_increment = max_level / max_bins;

	// Set upper and lower levels for each channel
	for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel)
	{
	int num_bins = num_levels[channel] - 1;
	lower_level[channel] = (max_level - (num_bins * min_level_increment)) / 2;
	upper_level[channel] = (max_level + (num_bins * min_level_increment)) / 2;
	}

	// Test pointer-based samples
	TestEvenNative<BACKEND, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleT, CounterT, LevelT, OffsetT>(
	max_level, entropy_reduction, num_levels, lower_level, upper_level, num_row_pixels, num_rows, row_stride_bytes);

	// Test iterator-based samples (CUB-only)
	TestEvenIterator<CUB, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleT, CounterT, LevelT, OffsetT>(
	max_level, entropy_reduction, num_levels, lower_level, upper_level, num_row_pixels, num_rows, row_stride_bytes);
	}



	/**
	* Test histogram-range
	*/
	template <
	Backend BACKEND,
	typename SampleT,
	int NUM_CHANNELS,
	int NUM_ACTIVE_CHANNELS,
	typename CounterT,
	typename LevelT,
	typename OffsetT>
	void TestRange(
	OffsetT num_row_pixels,
	OffsetT num_rows,
	OffsetT row_stride_bytes,
	int entropy_reduction,
	int num_levels[NUM_ACTIVE_CHANNELS],
	LevelT max_level,
	int max_num_levels)
	{
	// Find smallest level increment
	int max_bins = max_num_levels - 1;
	LevelT min_level_increment = max_level / max_bins;

	LevelT* levels[NUM_ACTIVE_CHANNELS];
	for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel)
	{
	levels[channel] = new LevelT[num_levels[channel]];

	int num_bins = num_levels[channel] - 1;
	LevelT lower_level = (max_level - (num_bins * min_level_increment)) / 2;

	for (int level = 0; level < num_levels[channel]; ++level)
	levels[channel][level] = lower_level + (level * min_level_increment);
	}

	TestRange<BACKEND, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleT, CounterT, LevelT, OffsetT>(
	max_level, entropy_reduction, num_levels, levels, num_row_pixels, num_rows, row_stride_bytes);

	for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel)
	delete[] levels[channel];

	}



	/**
	* Test different entrypoints
	*/
	template <
	typename SampleT,
	int NUM_CHANNELS,
	int NUM_ACTIVE_CHANNELS,
	typename CounterT,
	typename LevelT,
	typename OffsetT>
	void Test(
	OffsetT num_row_pixels,
	OffsetT num_rows,
	OffsetT row_stride_bytes,
	int entropy_reduction,
	int num_levels[NUM_ACTIVE_CHANNELS],
	LevelT max_level,
	int max_num_levels)
	{
	TestEven<CUB, SampleT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, CounterT, LevelT, OffsetT>(
	num_row_pixels, num_rows, row_stride_bytes, entropy_reduction, num_levels, max_level, max_num_levels);

	TestRange<CUB, SampleT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, CounterT, LevelT, OffsetT>(
	num_row_pixels, num_rows, row_stride_bytes, entropy_reduction, num_levels, max_level, max_num_levels);
	}


	/**
	* Test different number of levels
	*/
	template <
	typename SampleT,
	int NUM_CHANNELS,
	int NUM_ACTIVE_CHANNELS,
	typename CounterT,
	typename LevelT,
	typename OffsetT>
	void Test(
	OffsetT num_row_pixels,
	OffsetT num_rows,
	OffsetT row_stride_bytes,
	int entropy_reduction,
	LevelT max_level,
	int max_num_levels)
	{
	int num_levels[NUM_ACTIVE_CHANNELS];

	// Unnecessary testing
	// // All the same level
	// for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel)
	// {
	// num_levels[channel] = max_num_levels;
	// }
	// Test<SampleT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, CounterT, LevelT, OffsetT>(
	// num_row_pixels, num_rows, row_stride_bytes, entropy_reduction, num_levels, max_level, max_num_levels);

	// All different levels
	num_levels[0] = max_num_levels;
	for (int channel = 1; channel < NUM_ACTIVE_CHANNELS; ++channel)
	{
	num_levels[channel] = (num_levels[channel - 1] / 2) + 1;
	}
	Test<SampleT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, CounterT, LevelT, OffsetT>(
	num_row_pixels, num_rows, row_stride_bytes, entropy_reduction, num_levels, max_level, max_num_levels);
	}



	/**
	* Test different entropy-levels
	*/
	template <
	typename SampleT,
	int NUM_CHANNELS,
	int NUM_ACTIVE_CHANNELS,
	typename CounterT,
	typename LevelT,
	typename OffsetT>
	void Test(
	OffsetT num_row_pixels,
	OffsetT num_rows,
	OffsetT row_stride_bytes,
	LevelT max_level,
	int max_num_levels)
	{
	Test<SampleT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, CounterT, LevelT, OffsetT>(
	num_row_pixels, num_rows, row_stride_bytes, 0, max_level, max_num_levels);

	Test<SampleT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, CounterT, LevelT, OffsetT>(
	num_row_pixels, num_rows, row_stride_bytes, -1, max_level, max_num_levels);

	Test<SampleT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, CounterT, LevelT, OffsetT>(
	num_row_pixels, num_rows, row_stride_bytes, 5, max_level, max_num_levels);
	}


	/**
	* Test different row strides
	*/
	template <
	typename SampleT,
	int NUM_CHANNELS,
	int NUM_ACTIVE_CHANNELS,
	typename CounterT,
	typename LevelT,
	typename OffsetT>
	void Test(
	OffsetT num_row_pixels,
	OffsetT num_rows,
	LevelT max_level,
	int max_num_levels)
	{
	OffsetT row_stride_bytes = num_row_pixels * NUM_CHANNELS * sizeof(SampleT);

	// No padding
	Test<SampleT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, CounterT, LevelT, OffsetT>(
	num_row_pixels, num_rows, row_stride_bytes, max_level, max_num_levels);

	// 13 samples padding
	Test<SampleT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, CounterT, LevelT, OffsetT>(
	num_row_pixels, num_rows, row_stride_bytes + (13 * sizeof(SampleT)), max_level, max_num_levels);
	}


	/**
	* Test different problem sizes
	*/
	template <
	typename SampleT,
	int NUM_CHANNELS,
	int NUM_ACTIVE_CHANNELS,
	typename CounterT,
	typename LevelT,
	typename OffsetT>
	void Test(
	LevelT max_level,
	int max_num_levels)
	{
	// 0 row/col images
	Test<SampleT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, CounterT, LevelT, OffsetT>(
	OffsetT(1920), OffsetT(0), max_level, max_num_levels);
	Test<SampleT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, CounterT, LevelT, OffsetT>(
	OffsetT(0), OffsetT(0), max_level, max_num_levels);

	// 1080 image
	Test<SampleT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, CounterT, LevelT, OffsetT>(
	OffsetT(1920), OffsetT(1080), max_level, max_num_levels);

	// Sample different aspect ratios sizes
	for (OffsetT rows = 1; rows < 1000000; rows *= 1000)
	{
	for (OffsetT cols = 1; cols < (1000000 / rows); cols *= 1000)
	{
	Test<SampleT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, CounterT, LevelT, OffsetT>(
	cols, rows, max_level, max_num_levels);
	}
	}

	// Randomly select linear problem size between 1:10,000,000
	unsigned int max_int = (unsigned int) -1;
	for (int i = 0; i < 4; ++i)
	{
	unsigned int num_items;
	RandomBits(num_items);
	num_items = (unsigned int) ((double(num_items) * double(10000000)) / double(max_int));
	num_items = CUB_MAX(1, num_items);

	Test<SampleT, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, CounterT, LevelT, OffsetT>(
	OffsetT(num_items), 1, max_level, max_num_levels);
	}
	}



	/**
	* Test different channel interleavings (valid specialiation)
	*/
	template <
	typename SampleT,
	typename CounterT,
	typename LevelT,
	typename OffsetT>
	void TestChannels(
	LevelT max_level,
	int max_num_levels,
	Int2Type<true> /is_valid_tag/)
	{
	Test<SampleT, 1, 1, CounterT, LevelT, OffsetT>(max_level, max_num_levels);
	Test<SampleT, 4, 3, CounterT, LevelT, OffsetT>(max_level, max_num_levels);
	Test<SampleT, 3, 3, CounterT, LevelT, OffsetT>(max_level, max_num_levels);
	Test<SampleT, 4, 4, CounterT, LevelT, OffsetT>(max_level, max_num_levels);
	}


	/**
	* Test different channel interleavings (invalid specialiation)
	*/
	template <
	typename SampleT,
	typename CounterT,
	typename LevelT,
	typename OffsetT>
	void TestChannels(
	LevelT /max_level/,
	int /max_num_levels/,
	Int2Type<false> /is_valid_tag/)
	{}



	//---------------------------------------------------------------------
	// Main
	//---------------------------------------------------------------------




	/**
	* Main
	*/
	int main(int argc, char** argv)
	{
	int num_row_pixels = -1;
	int entropy_reduction = 0;
	int num_rows = 1;

	// Initialize command line
	CommandLineArgs args(argc, argv);
	g_verbose = args.CheckCmdLineFlag("v");
	g_verbose_input = args.CheckCmdLineFlag("v2");
	args.GetCmdLineArgument("n", num_row_pixels);

	int row_stride_pixels = num_row_pixels;

	args.GetCmdLineArgument("rows", num_rows);
	args.GetCmdLineArgument("stride", row_stride_pixels);
	args.GetCmdLineArgument("i", g_timing_iterations);
	args.GetCmdLineArgument("repeat", g_repeat);
	args.GetCmdLineArgument("entropy", entropy_reduction);
	#if defined(QUICK_TEST) \|\| defined(QUICKER_TEST)
	bool compare_npp = args.CheckCmdLineFlag("npp");
	#endif


	// Print usage
	if (args.CheckCmdLineFlag("help"))
	{
	printf("%s "
	"[--n=<pixels per row>] "
	"[--rows=<number of rows>] "
	"[--stride=<row stride in pixels>] "
	"[--i=<timing iterations>] "
	"[--device=<device-id>] "
	"[--repeat=<repetitions of entire test suite>]"
	"[--entropy=<entropy-reduction factor (default 0)>]"
	"[--v] "
	"[--cdp]"
	"[--npp]"
	"\n", argv[0]);
	exit(0);
	}

	// Initialize device
	CubDebugExit(args.DeviceInit());

	// Get ptx version
	int ptx_version = 0;
	CubDebugExit(PtxVersion(ptx_version));

	if (num_row_pixels < 0)
	{
	num_row_pixels = 1920 * 1080;
	row_stride_pixels = num_row_pixels;
	}

	#if defined(QUICKER_TEST)

	// Compile/run quick tests
	{
	// HistogramEven: unsigned char 256 bins
	typedef unsigned char SampleT;
	typedef int LevelT;

	LevelT max_level = 256;
	int num_levels[1] = {257};
	int row_stride_bytes = sizeof(SampleT) * row_stride_pixels * 1;

	TestEven<CUB, SampleT, 1, 1, int, LevelT, int>(num_row_pixels, num_rows, row_stride_bytes, entropy_reduction, num_levels, max_level, num_levels[0]);
	// The NPP path doesn't compile as of 2020-06:
	// No Dispatch<int, int, NPP> specialization defined.
	// if (compare_npp)
	// TestEven<NPP, SampleT, 1, 1, int, LevelT, int>(num_row_pixels, num_rows, row_stride_bytes, entropy_reduction, num_levels, max_level, num_levels[0]);
	}

	{
	// HistogramRange: signed char 256 bins
	typedef signed char SampleT;
	typedef int LevelT;

	LevelT max_level = 256;
	int num_levels[1] = {257};
	int row_stride_bytes = sizeof(SampleT) * row_stride_pixels * 1;

	TestRange<CUB, SampleT, 1, 1, int, LevelT, int>(num_row_pixels, num_rows, row_stride_bytes, entropy_reduction, num_levels, max_level, num_levels[0]);
	}



	#elif defined(QUICK_TEST)

	// Compile/run quick tests
	{
	// HistogramEven: unsigned char 256 bins
	typedef unsigned char SampleT;
	typedef int LevelT;

	LevelT max_level = 256;
	int num_levels[1] = {257};
	int row_stride_bytes = sizeof(SampleT) * row_stride_pixels * 1;

	TestEven<CUB, SampleT, 1, 1, int, LevelT, int>(num_row_pixels, num_rows, row_stride_bytes, entropy_reduction, num_levels, max_level, num_levels[0]);
	// The NPP path doesn't compile as of 2020-06:
	// No Dispatch<int, int, NPP> specialization defined.
	// if (compare_npp)
	// TestEven<NPP, SampleT, 1, 1, int, LevelT, int>(num_row_pixels, num_rows, row_stride_bytes, entropy_reduction, num_levels, max_level, num_levels[0]);
	}

	{
	// HistogramEven: 4/4 multichannel Unsigned char 256 bins
	typedef unsigned char SampleT;
	typedef int LevelT;

	LevelT max_level = 256;
	int num_levels[4] = {257, 257, 257, 257};
	int row_stride_bytes = sizeof(SampleT) * row_stride_pixels * 4;

	TestEven<CUB, SampleT, 4, 4, int, LevelT, int>(num_row_pixels, num_rows, row_stride_bytes, entropy_reduction, num_levels, max_level, num_levels[0]);
	}

	{
	// HistogramEven: 3/4 multichannel Unsigned char 256 bins
	typedef unsigned char SampleT;
	typedef int LevelT;

	LevelT max_level = 256;
	int num_levels[3] = {257, 257, 257};
	int row_stride_bytes = sizeof(SampleT) * row_stride_pixels * 4;

	TestEven<CUB, SampleT, 4, 3, int, LevelT, int>(num_row_pixels, num_rows, row_stride_bytes, entropy_reduction, num_levels, max_level, num_levels[0]);
	// The NPP path doesn't compile as of 2020-06:
	// No Dispatch<int, int, NPP> specialization defined.
	// if (compare_npp)
	// TestEven<NPP, SampleT, 4, 3, int, LevelT, int>(num_row_pixels, num_rows, row_stride_bytes, entropy_reduction, num_levels, max_level, num_levels[0]);
	}

	{
	// HistogramEven: short [0,1024] 256 bins
	typedef unsigned short SampleT;
	typedef unsigned short LevelT;

	LevelT max_level = 1024;
	int num_levels[1] = {257};
	int row_stride_bytes = sizeof(SampleT) * row_stride_pixels * 1;

	TestEven<CUB, SampleT, 1, 1, int, LevelT, int>(num_row_pixels, num_rows, row_stride_bytes, entropy_reduction, num_levels, max_level, num_levels[0]);
	}

	{
	// HistogramEven: float [0,1.0] 256 bins
	typedef float SampleT;
	typedef float LevelT;

	LevelT max_level = 1.0;
	int num_levels[1] = {257};
	int row_stride_bytes = sizeof(SampleT) * row_stride_pixels * 1;

	TestEven<CUB, SampleT, 1, 1, int, LevelT, int>(num_row_pixels, num_rows, row_stride_bytes, entropy_reduction, num_levels, max_level, num_levels[0]);
	}

	{
	// HistogramEven: 3/4 multichannel float [0,1.0] 256 bins
	typedef float SampleT;
	typedef float LevelT;

	LevelT max_level = 1.0;
	int num_levels[3] = {257, 257, 257};
	int row_stride_bytes = sizeof(SampleT) * row_stride_pixels * 4;

	TestEven<CUB, SampleT, 4, 3, int, LevelT, int>(num_row_pixels, num_rows, row_stride_bytes, entropy_reduction, num_levels, max_level, num_levels[0]);
	}

	{
	// HistogramRange: signed char 256 bins
	typedef signed char SampleT;
	typedef int LevelT;

	LevelT max_level = 256;
	int num_levels[1] = {257};
	int row_stride_bytes = sizeof(SampleT) * row_stride_pixels * 1;

	TestRange<CUB, SampleT, 1, 1, int, LevelT, int>(num_row_pixels, num_rows, row_stride_bytes, entropy_reduction, num_levels, max_level, num_levels[0]);
	}

	{
	// HistogramRange: 3/4 channel, unsigned char, varied bins (256, 128, 64)
	typedef unsigned char SampleT;
	typedef int LevelT;

	LevelT max_level = 256;
	int num_levels[3] = {257, 129, 65};
	int row_stride_bytes = sizeof(SampleT) * row_stride_pixels * 4;

	TestRange<CUB, SampleT, 4, 3, int, LevelT, int>(num_row_pixels, num_rows, row_stride_bytes, entropy_reduction, num_levels, max_level, num_levels[0]);
	}

	if (ptx_version > 120) // Don't check doubles on PTX120 or below because they're down-converted
	{
	// HistogramEven: double [0,1.0] 64 bins
	typedef double SampleT;
	typedef double LevelT;

	LevelT max_level = 1.0;
	int num_levels[1] = {65};
	int row_stride_bytes = sizeof(SampleT) * row_stride_pixels * 1;

	TestEven<CUB, SampleT, 1, 1, int, LevelT, int>(num_row_pixels, num_rows, row_stride_bytes, entropy_reduction, num_levels, max_level, num_levels[0]);
	}

	{
	// HistogramEven: short [0,1024] 512 bins
	typedef unsigned short SampleT;
	typedef unsigned short LevelT;

	LevelT max_level = 1024;
	int num_levels[1] = {513};
	int row_stride_bytes = sizeof(SampleT) * row_stride_pixels * 1;

	TestEven<CUB, SampleT, 1, 1, int, LevelT, int>(num_row_pixels, num_rows, row_stride_bytes, entropy_reduction, num_levels, max_level, num_levels[0]);
	}

	#else

	// Compile/run thorough tests
	for (int i = 0; i <= g_repeat; ++i)
	{
	TestChannels <unsigned char, int, int, int>(256, 256 + 1, Int2Type<true>());
	TestChannels <signed char, int, int, int>(256, 256 + 1, Int2Type<true>());
	TestChannels <unsigned short, int, int, int>(128, 128 + 1, Int2Type<true>());
	TestChannels <unsigned short, int, int, int>(8192, 8192 + 1, Int2Type<true>());
	TestChannels <float, int, float, int>(1.0, 256 + 1, Int2Type<true>());

	// Test down-conversion of size_t offsets to int
	TestChannels <unsigned char, int, int, long long>(256, 256 + 1, Int2Type<(sizeof(size_t) != sizeof(int))>());
	}

	#endif

	return 0;
	}