| /****************************************************************************** | |
| * 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; | |
| } | |