Spaces:
Runtime error
Runtime error
| /****************************************************************************** | |
| * 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. | |
| * | |
| ******************************************************************************/ | |
| /** | |
| * \file | |
| * cub::DeviceHistogram provides device-wide parallel operations for constructing histogram(s) from a sequence of samples data residing within device-accessible memory. | |
| */ | |
| #pragma once | |
| #include <stdio.h> | |
| #include <iterator> | |
| #include <limits> | |
| #include "../../agent/agent_histogram.cuh" | |
| #include "../../util_debug.cuh" | |
| #include "../../util_device.cuh" | |
| #include "../../thread/thread_search.cuh" | |
| #include "../../grid/grid_queue.cuh" | |
| #include "../../config.cuh" | |
| #include <thrust/system/cuda/detail/core/triple_chevron_launch.h> | |
| /// Optional outer namespace(s) | |
| CUB_NS_PREFIX | |
| /// CUB namespace | |
| namespace cub { | |
| /****************************************************************************** | |
| * Histogram kernel entry points | |
| *****************************************************************************/ | |
| /** | |
| * Histogram initialization kernel entry point | |
| */ | |
| template < | |
| int NUM_ACTIVE_CHANNELS, ///< Number of channels actively being histogrammed | |
| typename CounterT, ///< Integer type for counting sample occurrences per histogram bin | |
| typename OffsetT> ///< Signed integer type for global offsets | |
| __global__ void DeviceHistogramInitKernel( | |
| ArrayWrapper<int, NUM_ACTIVE_CHANNELS> num_output_bins_wrapper, ///< Number of output histogram bins per channel | |
| ArrayWrapper<CounterT*, NUM_ACTIVE_CHANNELS> d_output_histograms_wrapper, ///< Histogram counter data having logical dimensions <tt>CounterT[NUM_ACTIVE_CHANNELS][num_bins.array[CHANNEL]]</tt> | |
| GridQueue<int> tile_queue) ///< Drain queue descriptor for dynamically mapping tile data onto thread blocks | |
| { | |
| if ((threadIdx.x == 0) && (blockIdx.x == 0)) | |
| tile_queue.ResetDrain(); | |
| int output_bin = (blockIdx.x * blockDim.x) + threadIdx.x; | |
| #pragma unroll | |
| for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL) | |
| { | |
| if (output_bin < num_output_bins_wrapper.array[CHANNEL]) | |
| d_output_histograms_wrapper.array[CHANNEL][output_bin] = 0; | |
| } | |
| } | |
| /** | |
| * Histogram privatized sweep kernel entry point (multi-block). Computes privatized histograms, one per thread block. | |
| */ | |
| template < | |
| typename AgentHistogramPolicyT, ///< Parameterized AgentHistogramPolicy tuning policy type | |
| int PRIVATIZED_SMEM_BINS, ///< Maximum number of histogram bins per channel (e.g., up to 256) | |
| int NUM_CHANNELS, ///< Number of channels interleaved in the input data (may be greater than the number of channels being actively histogrammed) | |
| int NUM_ACTIVE_CHANNELS, ///< Number of channels actively being histogrammed | |
| typename SampleIteratorT, ///< The input iterator type. \iterator. | |
| typename CounterT, ///< Integer type for counting sample occurrences per histogram bin | |
| typename PrivatizedDecodeOpT, ///< The transform operator type for determining privatized counter indices from samples, one for each channel | |
| typename OutputDecodeOpT, ///< The transform operator type for determining output bin-ids from privatized counter indices, one for each channel | |
| typename OffsetT> ///< Signed integer type for global offsets | |
| __launch_bounds__ (int(AgentHistogramPolicyT::BLOCK_THREADS)) | |
| __global__ void DeviceHistogramSweepKernel( | |
| SampleIteratorT d_samples, ///< Input data to reduce | |
| ArrayWrapper<int, NUM_ACTIVE_CHANNELS> num_output_bins_wrapper, ///< The number bins per final output histogram | |
| ArrayWrapper<int, NUM_ACTIVE_CHANNELS> num_privatized_bins_wrapper, ///< The number bins per privatized histogram | |
| ArrayWrapper<CounterT*, NUM_ACTIVE_CHANNELS> d_output_histograms_wrapper, ///< Reference to final output histograms | |
| ArrayWrapper<CounterT*, NUM_ACTIVE_CHANNELS> d_privatized_histograms_wrapper, ///< Reference to privatized histograms | |
| ArrayWrapper<OutputDecodeOpT, NUM_ACTIVE_CHANNELS> output_decode_op_wrapper, ///< The transform operator for determining output bin-ids from privatized counter indices, one for each channel | |
| ArrayWrapper<PrivatizedDecodeOpT, NUM_ACTIVE_CHANNELS> privatized_decode_op_wrapper, ///< The transform operator for determining privatized counter indices from samples, one for each channel | |
| OffsetT num_row_pixels, ///< The number of multi-channel pixels per row in the region of interest | |
| OffsetT num_rows, ///< The number of rows in the region of interest | |
| OffsetT row_stride_samples, ///< The number of samples between starts of consecutive rows in the region of interest | |
| int tiles_per_row, ///< Number of image tiles per row | |
| GridQueue<int> tile_queue) ///< Drain queue descriptor for dynamically mapping tile data onto thread blocks | |
| { | |
| // Thread block type for compositing input tiles | |
| typedef AgentHistogram< | |
| AgentHistogramPolicyT, | |
| PRIVATIZED_SMEM_BINS, | |
| NUM_CHANNELS, | |
| NUM_ACTIVE_CHANNELS, | |
| SampleIteratorT, | |
| CounterT, | |
| PrivatizedDecodeOpT, | |
| OutputDecodeOpT, | |
| OffsetT> | |
| AgentHistogramT; | |
| // Shared memory for AgentHistogram | |
| __shared__ typename AgentHistogramT::TempStorage temp_storage; | |
| AgentHistogramT agent( | |
| temp_storage, | |
| d_samples, | |
| num_output_bins_wrapper.array, | |
| num_privatized_bins_wrapper.array, | |
| d_output_histograms_wrapper.array, | |
| d_privatized_histograms_wrapper.array, | |
| output_decode_op_wrapper.array, | |
| privatized_decode_op_wrapper.array); | |
| // Initialize counters | |
| agent.InitBinCounters(); | |
| // Consume input tiles | |
| agent.ConsumeTiles( | |
| num_row_pixels, | |
| num_rows, | |
| row_stride_samples, | |
| tiles_per_row, | |
| tile_queue); | |
| // Store output to global (if necessary) | |
| agent.StoreOutput(); | |
| } | |
| /****************************************************************************** | |
| * Dispatch | |
| ******************************************************************************/ | |
| /** | |
| * Utility class for dispatching the appropriately-tuned kernels for DeviceHistogram | |
| */ | |
| template < | |
| int NUM_CHANNELS, ///< Number of channels interleaved in the input data (may be greater than the number of channels being actively histogrammed) | |
| int NUM_ACTIVE_CHANNELS, ///< Number of channels actively being histogrammed | |
| typename SampleIteratorT, ///< Random-access input iterator type for reading input items \iterator | |
| typename CounterT, ///< Integer type for counting sample occurrences per histogram bin | |
| typename LevelT, ///< Type for specifying bin level boundaries | |
| typename OffsetT> ///< Signed integer type for global offsets | |
| struct DipatchHistogram | |
| { | |
| //--------------------------------------------------------------------- | |
| // Types and constants | |
| //--------------------------------------------------------------------- | |
| /// The sample value type of the input iterator | |
| typedef typename std::iterator_traits<SampleIteratorT>::value_type SampleT; | |
| enum | |
| { | |
| // Maximum number of bins per channel for which we will use a privatized smem strategy | |
| MAX_PRIVATIZED_SMEM_BINS = 256 | |
| }; | |
| //--------------------------------------------------------------------- | |
| // Transform functors for converting samples to bin-ids | |
| //--------------------------------------------------------------------- | |
| // Searches for bin given a list of bin-boundary levels | |
| template <typename LevelIteratorT> | |
| struct SearchTransform | |
| { | |
| LevelIteratorT d_levels; // Pointer to levels array | |
| int num_output_levels; // Number of levels in array | |
| // Initializer | |
| __host__ __device__ __forceinline__ void Init( | |
| LevelIteratorT d_levels, // Pointer to levels array | |
| int num_output_levels) // Number of levels in array | |
| { | |
| this->d_levels = d_levels; | |
| this->num_output_levels = num_output_levels; | |
| } | |
| // Method for converting samples to bin-ids | |
| template <CacheLoadModifier LOAD_MODIFIER, typename _SampleT> | |
| __host__ __device__ __forceinline__ void BinSelect(_SampleT sample, int &bin, bool valid) | |
| { | |
| /// Level iterator wrapper type | |
| typedef typename If<IsPointer<LevelIteratorT>::VALUE, | |
| CacheModifiedInputIterator<LOAD_MODIFIER, LevelT, OffsetT>, // Wrap the native input pointer with CacheModifiedInputIterator | |
| LevelIteratorT>::Type // Directly use the supplied input iterator type | |
| WrappedLevelIteratorT; | |
| WrappedLevelIteratorT wrapped_levels(d_levels); | |
| int num_bins = num_output_levels - 1; | |
| if (valid) | |
| { | |
| bin = UpperBound(wrapped_levels, num_output_levels, (LevelT) sample) - 1; | |
| if (bin >= num_bins) | |
| bin = -1; | |
| } | |
| } | |
| }; | |
| // Scales samples to evenly-spaced bins | |
| struct ScaleTransform | |
| { | |
| int num_bins; // Number of levels in array | |
| LevelT max; // Max sample level (exclusive) | |
| LevelT min; // Min sample level (inclusive) | |
| LevelT scale; // Bin scaling factor | |
| // Initializer | |
| template <typename _LevelT> | |
| __host__ __device__ __forceinline__ void Init( | |
| int num_output_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_bins = num_output_levels - 1; | |
| this->max = max; | |
| this->min = min; | |
| this->scale = scale; | |
| } | |
| // Initializer (float specialization) | |
| __host__ __device__ __forceinline__ void Init( | |
| int num_output_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_bins = num_output_levels - 1; | |
| this->max = max; | |
| this->min = min; | |
| this->scale = float(1.0) / scale; | |
| } | |
| // Initializer (double specialization) | |
| __host__ __device__ __forceinline__ void Init( | |
| int num_output_levels, // Number of levels in array | |
| double max, // Max sample level (exclusive) | |
| double min, // Min sample level (inclusive) | |
| double scale) // Bin scaling factor | |
| { | |
| this->num_bins = num_output_levels - 1; | |
| this->max = max; | |
| this->min = min; | |
| this->scale = double(1.0) / scale; | |
| } | |
| // Method for converting samples to bin-ids | |
| template <CacheLoadModifier LOAD_MODIFIER, typename _SampleT> | |
| __host__ __device__ __forceinline__ void BinSelect(_SampleT sample, int &bin, bool valid) | |
| { | |
| LevelT level_sample = (LevelT) sample; | |
| if (valid && (level_sample >= min) && (level_sample < max)) | |
| bin = (int) ((level_sample - min) / scale); | |
| } | |
| // Method for converting samples to bin-ids (float specialization) | |
| template <CacheLoadModifier LOAD_MODIFIER> | |
| __host__ __device__ __forceinline__ void BinSelect(float sample, int &bin, bool valid) | |
| { | |
| LevelT level_sample = (LevelT) sample; | |
| if (valid && (level_sample >= min) && (level_sample < max)) | |
| bin = (int) ((level_sample - min) * scale); | |
| } | |
| // Method for converting samples to bin-ids (double specialization) | |
| template <CacheLoadModifier LOAD_MODIFIER> | |
| __host__ __device__ __forceinline__ void BinSelect(double sample, int &bin, bool valid) | |
| { | |
| LevelT level_sample = (LevelT) sample; | |
| if (valid && (level_sample >= min) && (level_sample < max)) | |
| bin = (int) ((level_sample - min) * scale); | |
| } | |
| }; | |
| // Pass-through bin transform operator | |
| struct PassThruTransform | |
| { | |
| // Method for converting samples to bin-ids | |
| template <CacheLoadModifier LOAD_MODIFIER, typename _SampleT> | |
| __host__ __device__ __forceinline__ void BinSelect(_SampleT sample, int &bin, bool valid) | |
| { | |
| if (valid) | |
| bin = (int) sample; | |
| } | |
| }; | |
| //--------------------------------------------------------------------- | |
| // Tuning policies | |
| //--------------------------------------------------------------------- | |
| template <int NOMINAL_ITEMS_PER_THREAD> | |
| struct TScale | |
| { | |
| enum | |
| { | |
| V_SCALE = (sizeof(SampleT) + sizeof(int) - 1) / sizeof(int), | |
| VALUE = CUB_MAX((NOMINAL_ITEMS_PER_THREAD / NUM_ACTIVE_CHANNELS / V_SCALE), 1) | |
| }; | |
| }; | |
| /// SM11 | |
| struct Policy110 | |
| { | |
| // HistogramSweepPolicy | |
| typedef AgentHistogramPolicy< | |
| 512, | |
| (NUM_CHANNELS == 1) ? 8 : 2, | |
| BLOCK_LOAD_DIRECT, | |
| LOAD_DEFAULT, | |
| true, | |
| GMEM, | |
| false> | |
| HistogramSweepPolicy; | |
| }; | |
| /// SM20 | |
| struct Policy200 | |
| { | |
| // HistogramSweepPolicy | |
| typedef AgentHistogramPolicy< | |
| (NUM_CHANNELS == 1) ? 256 : 128, | |
| (NUM_CHANNELS == 1) ? 8 : 3, | |
| (NUM_CHANNELS == 1) ? BLOCK_LOAD_DIRECT : BLOCK_LOAD_WARP_TRANSPOSE, | |
| LOAD_DEFAULT, | |
| true, | |
| SMEM, | |
| false> | |
| HistogramSweepPolicy; | |
| }; | |
| /// SM30 | |
| struct Policy300 | |
| { | |
| // HistogramSweepPolicy | |
| typedef AgentHistogramPolicy< | |
| 512, | |
| (NUM_CHANNELS == 1) ? 8 : 2, | |
| BLOCK_LOAD_DIRECT, | |
| LOAD_DEFAULT, | |
| true, | |
| GMEM, | |
| false> | |
| HistogramSweepPolicy; | |
| }; | |
| /// SM35 | |
| struct Policy350 | |
| { | |
| // HistogramSweepPolicy | |
| typedef AgentHistogramPolicy< | |
| 128, | |
| TScale<8>::VALUE, | |
| BLOCK_LOAD_DIRECT, | |
| LOAD_LDG, | |
| true, | |
| BLEND, | |
| true> | |
| HistogramSweepPolicy; | |
| }; | |
| /// SM50 | |
| struct Policy500 | |
| { | |
| // HistogramSweepPolicy | |
| typedef AgentHistogramPolicy< | |
| 384, | |
| TScale<16>::VALUE, | |
| BLOCK_LOAD_DIRECT, | |
| LOAD_LDG, | |
| true, | |
| SMEM, | |
| false> | |
| HistogramSweepPolicy; | |
| }; | |
| //--------------------------------------------------------------------- | |
| // Tuning policies of current PTX compiler pass | |
| //--------------------------------------------------------------------- | |
| #if (CUB_PTX_ARCH >= 500) | |
| typedef Policy500 PtxPolicy; | |
| #elif (CUB_PTX_ARCH >= 350) | |
| typedef Policy350 PtxPolicy; | |
| #elif (CUB_PTX_ARCH >= 300) | |
| typedef Policy300 PtxPolicy; | |
| #elif (CUB_PTX_ARCH >= 200) | |
| typedef Policy200 PtxPolicy; | |
| #else | |
| typedef Policy110 PtxPolicy; | |
| #endif | |
| // "Opaque" policies (whose parameterizations aren't reflected in the type signature) | |
| struct PtxHistogramSweepPolicy : PtxPolicy::HistogramSweepPolicy {}; | |
| //--------------------------------------------------------------------- | |
| // Utilities | |
| //--------------------------------------------------------------------- | |
| /** | |
| * Initialize kernel dispatch configurations with the policies corresponding to the PTX assembly we will use | |
| */ | |
| template <typename KernelConfig> | |
| CUB_RUNTIME_FUNCTION __forceinline__ | |
| static cudaError_t InitConfigs( | |
| int ptx_version, | |
| KernelConfig &histogram_sweep_config) | |
| { | |
| cudaError_t result = cudaErrorNotSupported; | |
| if (CUB_IS_DEVICE_CODE) | |
| { | |
| #if CUB_INCLUDE_DEVICE_CODE | |
| // We're on the device, so initialize the kernel dispatch configurations with the current PTX policy | |
| result = histogram_sweep_config.template Init<PtxHistogramSweepPolicy>(); | |
| #endif | |
| } | |
| else | |
| { | |
| #if CUB_INCLUDE_HOST_CODE | |
| // We're on the host, so lookup and initialize the kernel dispatch configurations with the policies that match the device's PTX version | |
| if (ptx_version >= 500) | |
| { | |
| result = histogram_sweep_config.template Init<typename Policy500::HistogramSweepPolicy>(); | |
| } | |
| else if (ptx_version >= 350) | |
| { | |
| result = histogram_sweep_config.template Init<typename Policy350::HistogramSweepPolicy>(); | |
| } | |
| else if (ptx_version >= 300) | |
| { | |
| result = histogram_sweep_config.template Init<typename Policy300::HistogramSweepPolicy>(); | |
| } | |
| else if (ptx_version >= 200) | |
| { | |
| result = histogram_sweep_config.template Init<typename Policy200::HistogramSweepPolicy>(); | |
| } | |
| else | |
| { | |
| result = histogram_sweep_config.template Init<typename Policy110::HistogramSweepPolicy>(); | |
| } | |
| #endif | |
| } | |
| return result; | |
| } | |
| /** | |
| * Kernel kernel dispatch configuration | |
| */ | |
| struct KernelConfig | |
| { | |
| int block_threads; | |
| int pixels_per_thread; | |
| template <typename BlockPolicy> | |
| CUB_RUNTIME_FUNCTION __forceinline__ | |
| cudaError_t Init() | |
| { | |
| block_threads = BlockPolicy::BLOCK_THREADS; | |
| pixels_per_thread = BlockPolicy::PIXELS_PER_THREAD; | |
| return cudaSuccess; | |
| } | |
| }; | |
| //--------------------------------------------------------------------- | |
| // Dispatch entrypoints | |
| //--------------------------------------------------------------------- | |
| /** | |
| * Privatization-based dispatch routine | |
| */ | |
| template < | |
| typename PrivatizedDecodeOpT, ///< The transform operator type for determining privatized counter indices from samples, one for each channel | |
| typename OutputDecodeOpT, ///< The transform operator type for determining output bin-ids from privatized counter indices, one for each channel | |
| typename DeviceHistogramInitKernelT, ///< Function type of cub::DeviceHistogramInitKernel | |
| typename DeviceHistogramSweepKernelT> ///< Function type of cub::DeviceHistogramSweepKernel | |
| CUB_RUNTIME_FUNCTION __forceinline__ | |
| static cudaError_t PrivatizedDispatch( | |
| void* d_temp_storage, ///< [in] %Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to \p temp_storage_bytes and no work is done. | |
| size_t& temp_storage_bytes, ///< [in,out] Reference to size in bytes of \p d_temp_storage allocation | |
| SampleIteratorT d_samples, ///< [in] The pointer to the input sequence of sample items. 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_output_histograms[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_output_levels[i]</tt> - 1. | |
| int num_privatized_levels[NUM_ACTIVE_CHANNELS], ///< [in] The number of bin level boundaries for delineating histogram samples in each active channel. Implies that the number of bins for channel<sub><em>i</em></sub> is <tt>num_output_levels[i]</tt> - 1. | |
| PrivatizedDecodeOpT privatized_decode_op[NUM_ACTIVE_CHANNELS], ///< [in] Transform operators for determining bin-ids from samples, one for each channel | |
| int num_output_levels[NUM_ACTIVE_CHANNELS], ///< [in] The number of bin level boundaries for delineating histogram samples in each active channel. Implies that the number of bins for channel<sub><em>i</em></sub> is <tt>num_output_levels[i]</tt> - 1. | |
| OutputDecodeOpT output_decode_op[NUM_ACTIVE_CHANNELS], ///< [in] Transform operators for determining bin-ids from samples, one for each channel | |
| int max_num_output_bins, ///< [in] Maximum number of output bins in any 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_samples, ///< [in] The number of samples between starts of consecutive rows in the region of interest | |
| DeviceHistogramInitKernelT histogram_init_kernel, ///< [in] Kernel function pointer to parameterization of cub::DeviceHistogramInitKernel | |
| DeviceHistogramSweepKernelT histogram_sweep_kernel, ///< [in] Kernel function pointer to parameterization of cub::DeviceHistogramSweepKernel | |
| KernelConfig histogram_sweep_config, ///< [in] Dispatch parameters that match the policy that \p histogram_sweep_kernel was compiled for | |
| cudaStream_t stream, ///< [in] CUDA stream to launch kernels within. Default is stream<sub>0</sub>. | |
| bool debug_synchronous) ///< [in] Whether or not to synchronize the stream after every kernel launch to check for errors. May cause significant slowdown. Default is \p false. | |
| { | |
| #ifndef CUB_RUNTIME_ENABLED | |
| // Kernel launch not supported from this device | |
| return CubDebug(cudaErrorNotSupported); | |
| #else | |
| cudaError error = cudaSuccess; | |
| do | |
| { | |
| // Get device ordinal | |
| int device_ordinal; | |
| if (CubDebug(error = cudaGetDevice(&device_ordinal))) break; | |
| // Get SM count | |
| int sm_count; | |
| if (CubDebug(error = cudaDeviceGetAttribute (&sm_count, cudaDevAttrMultiProcessorCount, device_ordinal))) break; | |
| // Get SM occupancy for histogram_sweep_kernel | |
| int histogram_sweep_sm_occupancy; | |
| if (CubDebug(error = MaxSmOccupancy( | |
| histogram_sweep_sm_occupancy, | |
| histogram_sweep_kernel, | |
| histogram_sweep_config.block_threads))) break; | |
| // Get device occupancy for histogram_sweep_kernel | |
| int histogram_sweep_occupancy = histogram_sweep_sm_occupancy * sm_count; | |
| if (num_row_pixels * NUM_CHANNELS == row_stride_samples) | |
| { | |
| // Treat as a single linear array of samples | |
| num_row_pixels *= num_rows; | |
| num_rows = 1; | |
| row_stride_samples = num_row_pixels * NUM_CHANNELS; | |
| } | |
| // Get grid dimensions, trying to keep total blocks ~histogram_sweep_occupancy | |
| int pixels_per_tile = histogram_sweep_config.block_threads * histogram_sweep_config.pixels_per_thread; | |
| int tiles_per_row = int(num_row_pixels + pixels_per_tile - 1) / pixels_per_tile; | |
| int blocks_per_row = CUB_MIN(histogram_sweep_occupancy, tiles_per_row); | |
| int blocks_per_col = (blocks_per_row > 0) ? | |
| int(CUB_MIN(histogram_sweep_occupancy / blocks_per_row, num_rows)) : | |
| 0; | |
| int num_thread_blocks = blocks_per_row * blocks_per_col; | |
| dim3 sweep_grid_dims; | |
| sweep_grid_dims.x = (unsigned int) blocks_per_row; | |
| sweep_grid_dims.y = (unsigned int) blocks_per_col; | |
| sweep_grid_dims.z = 1; | |
| // Temporary storage allocation requirements | |
| const int NUM_ALLOCATIONS = NUM_ACTIVE_CHANNELS + 1; | |
| void* allocations[NUM_ALLOCATIONS] = {}; | |
| size_t allocation_sizes[NUM_ALLOCATIONS]; | |
| for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL) | |
| allocation_sizes[CHANNEL] = size_t(num_thread_blocks) * (num_privatized_levels[CHANNEL] - 1) * sizeof(CounterT); | |
| allocation_sizes[NUM_ALLOCATIONS - 1] = GridQueue<int>::AllocationSize(); | |
| // Alias the temporary allocations from the single storage blob (or compute the necessary size of the blob) | |
| if (CubDebug(error = AliasTemporaries(d_temp_storage, temp_storage_bytes, allocations, allocation_sizes))) break; | |
| if (d_temp_storage == NULL) | |
| { | |
| // Return if the caller is simply requesting the size of the storage allocation | |
| break; | |
| } | |
| // Construct the grid queue descriptor | |
| GridQueue<int> tile_queue(allocations[NUM_ALLOCATIONS - 1]); | |
| // Setup array wrapper for histogram channel output (because we can't pass static arrays as kernel parameters) | |
| ArrayWrapper<CounterT*, NUM_ACTIVE_CHANNELS> d_output_histograms_wrapper; | |
| for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL) | |
| d_output_histograms_wrapper.array[CHANNEL] = d_output_histograms[CHANNEL]; | |
| // Setup array wrapper for privatized per-block histogram channel output (because we can't pass static arrays as kernel parameters) | |
| ArrayWrapper<CounterT*, NUM_ACTIVE_CHANNELS> d_privatized_histograms_wrapper; | |
| for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL) | |
| d_privatized_histograms_wrapper.array[CHANNEL] = (CounterT*) allocations[CHANNEL]; | |
| // Setup array wrapper for sweep bin transforms (because we can't pass static arrays as kernel parameters) | |
| ArrayWrapper<PrivatizedDecodeOpT, NUM_ACTIVE_CHANNELS> privatized_decode_op_wrapper; | |
| for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL) | |
| privatized_decode_op_wrapper.array[CHANNEL] = privatized_decode_op[CHANNEL]; | |
| // Setup array wrapper for aggregation bin transforms (because we can't pass static arrays as kernel parameters) | |
| ArrayWrapper<OutputDecodeOpT, NUM_ACTIVE_CHANNELS> output_decode_op_wrapper; | |
| for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL) | |
| output_decode_op_wrapper.array[CHANNEL] = output_decode_op[CHANNEL]; | |
| // Setup array wrapper for num privatized bins (because we can't pass static arrays as kernel parameters) | |
| ArrayWrapper<int, NUM_ACTIVE_CHANNELS> num_privatized_bins_wrapper; | |
| for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL) | |
| num_privatized_bins_wrapper.array[CHANNEL] = num_privatized_levels[CHANNEL] - 1; | |
| // Setup array wrapper for num output bins (because we can't pass static arrays as kernel parameters) | |
| ArrayWrapper<int, NUM_ACTIVE_CHANNELS> num_output_bins_wrapper; | |
| for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL) | |
| num_output_bins_wrapper.array[CHANNEL] = num_output_levels[CHANNEL] - 1; | |
| int histogram_init_block_threads = 256; | |
| int histogram_init_grid_dims = (max_num_output_bins + histogram_init_block_threads - 1) / histogram_init_block_threads; | |
| // Log DeviceHistogramInitKernel configuration | |
| if (debug_synchronous) _CubLog("Invoking DeviceHistogramInitKernel<<<%d, %d, 0, %lld>>>()\n", | |
| histogram_init_grid_dims, histogram_init_block_threads, (long long) stream); | |
| // Invoke histogram_init_kernel | |
| thrust::cuda_cub::launcher::triple_chevron( | |
| histogram_init_grid_dims, histogram_init_block_threads, 0, | |
| stream | |
| ).doit(histogram_init_kernel, | |
| num_output_bins_wrapper, | |
| d_output_histograms_wrapper, | |
| tile_queue); | |
| // Return if empty problem | |
| if ((blocks_per_row == 0) || (blocks_per_col == 0)) | |
| break; | |
| // Log histogram_sweep_kernel configuration | |
| if (debug_synchronous) _CubLog("Invoking histogram_sweep_kernel<<<{%d, %d, %d}, %d, 0, %lld>>>(), %d pixels per thread, %d SM occupancy\n", | |
| sweep_grid_dims.x, sweep_grid_dims.y, sweep_grid_dims.z, | |
| histogram_sweep_config.block_threads, (long long) stream, histogram_sweep_config.pixels_per_thread, histogram_sweep_sm_occupancy); | |
| // Invoke histogram_sweep_kernel | |
| thrust::cuda_cub::launcher::triple_chevron( | |
| sweep_grid_dims, histogram_sweep_config.block_threads, 0, stream | |
| ).doit(histogram_sweep_kernel, | |
| d_samples, | |
| num_output_bins_wrapper, | |
| num_privatized_bins_wrapper, | |
| d_output_histograms_wrapper, | |
| d_privatized_histograms_wrapper, | |
| output_decode_op_wrapper, | |
| privatized_decode_op_wrapper, | |
| num_row_pixels, | |
| num_rows, | |
| row_stride_samples, | |
| tiles_per_row, | |
| tile_queue); | |
| // Check for failure to launch | |
| if (CubDebug(error = cudaPeekAtLastError())) break; | |
| // Sync the stream if specified to flush runtime errors | |
| if (debug_synchronous && (CubDebug(error = SyncStream(stream)))) break; | |
| } | |
| while (0); | |
| return error; | |
| #endif // CUB_RUNTIME_ENABLED | |
| } | |
| /** | |
| * Dispatch routine for HistogramRange, specialized for sample types larger than 8bit | |
| */ | |
| CUB_RUNTIME_FUNCTION | |
| static cudaError_t DispatchRange( | |
| void* d_temp_storage, ///< [in] %Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to \p temp_storage_bytes and no work is done. | |
| size_t& temp_storage_bytes, ///< [in,out] Reference to size in bytes of \p d_temp_storage allocation | |
| 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_output_histograms[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_output_levels[i]</tt> - 1. | |
| int num_output_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_output_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_samples, ///< [in] The number of samples between starts of consecutive rows in the region of interest | |
| cudaStream_t stream, ///< [in] CUDA stream to launch kernels within. Default is stream<sub>0</sub>. | |
| bool debug_synchronous, ///< [in] Whether or not to synchronize the stream after every kernel launch to check for errors. May cause significant slowdown. Default is \p false. | |
| Int2Type<false> /*is_byte_sample*/) ///< [in] Marker type indicating whether or not SampleT is a 8b type | |
| { | |
| cudaError error = cudaSuccess; | |
| do | |
| { | |
| // Get PTX version | |
| int ptx_version = 0; | |
| if (CubDebug(error = PtxVersion(ptx_version))) break; | |
| // Get kernel dispatch configurations | |
| KernelConfig histogram_sweep_config; | |
| if (CubDebug(error = InitConfigs(ptx_version, histogram_sweep_config))) | |
| break; | |
| // Use the search transform op for converting samples to privatized bins | |
| typedef SearchTransform<LevelT*> PrivatizedDecodeOpT; | |
| // Use the pass-thru transform op for converting privatized bins to output bins | |
| typedef PassThruTransform OutputDecodeOpT; | |
| PrivatizedDecodeOpT privatized_decode_op[NUM_ACTIVE_CHANNELS]; | |
| OutputDecodeOpT output_decode_op[NUM_ACTIVE_CHANNELS]; | |
| int max_levels = num_output_levels[0]; | |
| for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel) | |
| { | |
| privatized_decode_op[channel].Init(d_levels[channel], num_output_levels[channel]); | |
| if (num_output_levels[channel] > max_levels) | |
| max_levels = num_output_levels[channel]; | |
| } | |
| int max_num_output_bins = max_levels - 1; | |
| // Dispatch | |
| if (max_num_output_bins > MAX_PRIVATIZED_SMEM_BINS) | |
| { | |
| // Too many bins to keep in shared memory. | |
| const int PRIVATIZED_SMEM_BINS = 0; | |
| if (CubDebug(error = PrivatizedDispatch( | |
| d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_output_histograms, | |
| num_output_levels, | |
| privatized_decode_op, | |
| num_output_levels, | |
| output_decode_op, | |
| max_num_output_bins, | |
| num_row_pixels, | |
| num_rows, | |
| row_stride_samples, | |
| DeviceHistogramInitKernel<NUM_ACTIVE_CHANNELS, CounterT, OffsetT>, | |
| DeviceHistogramSweepKernel<PtxHistogramSweepPolicy, PRIVATIZED_SMEM_BINS, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleIteratorT, CounterT, PrivatizedDecodeOpT, OutputDecodeOpT, OffsetT>, | |
| histogram_sweep_config, | |
| stream, | |
| debug_synchronous))) break; | |
| } | |
| else | |
| { | |
| // Dispatch shared-privatized approach | |
| const int PRIVATIZED_SMEM_BINS = MAX_PRIVATIZED_SMEM_BINS; | |
| if (CubDebug(error = PrivatizedDispatch( | |
| d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_output_histograms, | |
| num_output_levels, | |
| privatized_decode_op, | |
| num_output_levels, | |
| output_decode_op, | |
| max_num_output_bins, | |
| num_row_pixels, | |
| num_rows, | |
| row_stride_samples, | |
| DeviceHistogramInitKernel<NUM_ACTIVE_CHANNELS, CounterT, OffsetT>, | |
| DeviceHistogramSweepKernel<PtxHistogramSweepPolicy, PRIVATIZED_SMEM_BINS, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleIteratorT, CounterT, PrivatizedDecodeOpT, OutputDecodeOpT, OffsetT>, | |
| histogram_sweep_config, | |
| stream, | |
| debug_synchronous))) break; | |
| } | |
| } while (0); | |
| return error; | |
| } | |
| /** | |
| * Dispatch routine for HistogramRange, specialized for 8-bit sample types (computes 256-bin privatized histograms and then reduces to user-specified levels) | |
| */ | |
| CUB_RUNTIME_FUNCTION | |
| static cudaError_t DispatchRange( | |
| void* d_temp_storage, ///< [in] %Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to \p temp_storage_bytes and no work is done. | |
| size_t& temp_storage_bytes, ///< [in,out] Reference to size in bytes of \p d_temp_storage allocation | |
| 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_output_histograms[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_output_levels[i]</tt> - 1. | |
| int num_output_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_output_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_samples, ///< [in] The number of samples between starts of consecutive rows in the region of interest | |
| cudaStream_t stream, ///< [in] CUDA stream to launch kernels within. Default is stream<sub>0</sub>. | |
| bool debug_synchronous, ///< [in] Whether or not to synchronize the stream after every kernel launch to check for errors. May cause significant slowdown. Default is \p false. | |
| Int2Type<true> /*is_byte_sample*/) ///< [in] Marker type indicating whether or not SampleT is a 8b type | |
| { | |
| cudaError error = cudaSuccess; | |
| do | |
| { | |
| // Get PTX version | |
| int ptx_version = 0; | |
| if (CubDebug(error = PtxVersion(ptx_version))) break; | |
| // Get kernel dispatch configurations | |
| KernelConfig histogram_sweep_config; | |
| if (CubDebug(error = InitConfigs(ptx_version, histogram_sweep_config))) | |
| break; | |
| // Use the pass-thru transform op for converting samples to privatized bins | |
| typedef PassThruTransform PrivatizedDecodeOpT; | |
| // Use the search transform op for converting privatized bins to output bins | |
| typedef SearchTransform<LevelT*> OutputDecodeOpT; | |
| int num_privatized_levels[NUM_ACTIVE_CHANNELS]; | |
| PrivatizedDecodeOpT privatized_decode_op[NUM_ACTIVE_CHANNELS]; | |
| OutputDecodeOpT output_decode_op[NUM_ACTIVE_CHANNELS]; | |
| int max_levels = num_output_levels[0]; // Maximum number of levels in any channel | |
| for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel) | |
| { | |
| num_privatized_levels[channel] = 257; | |
| output_decode_op[channel].Init(d_levels[channel], num_output_levels[channel]); | |
| if (num_output_levels[channel] > max_levels) | |
| max_levels = num_output_levels[channel]; | |
| } | |
| int max_num_output_bins = max_levels - 1; | |
| const int PRIVATIZED_SMEM_BINS = 256; | |
| if (CubDebug(error = PrivatizedDispatch( | |
| d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_output_histograms, | |
| num_privatized_levels, | |
| privatized_decode_op, | |
| num_output_levels, | |
| output_decode_op, | |
| max_num_output_bins, | |
| num_row_pixels, | |
| num_rows, | |
| row_stride_samples, | |
| DeviceHistogramInitKernel<NUM_ACTIVE_CHANNELS, CounterT, OffsetT>, | |
| DeviceHistogramSweepKernel<PtxHistogramSweepPolicy, PRIVATIZED_SMEM_BINS, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleIteratorT, CounterT, PrivatizedDecodeOpT, OutputDecodeOpT, OffsetT>, | |
| histogram_sweep_config, | |
| stream, | |
| debug_synchronous))) break; | |
| } while (0); | |
| return error; | |
| } | |
| /** | |
| * Dispatch routine for HistogramEven, specialized for sample types larger than 8-bit | |
| */ | |
| CUB_RUNTIME_FUNCTION __forceinline__ | |
| static cudaError_t DispatchEven( | |
| void* d_temp_storage, ///< [in] %Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to \p temp_storage_bytes and no work is done. | |
| size_t& temp_storage_bytes, ///< [in,out] Reference to size in bytes of \p d_temp_storage allocation | |
| SampleIteratorT d_samples, ///< [in] The pointer to the input sequence of sample items. 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_output_histograms[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_output_levels[i]</tt> - 1. | |
| int num_output_levels[NUM_ACTIVE_CHANNELS], ///< [in] The number of bin level boundaries for delineating histogram samples in each active channel. Implies that the number of bins for channel<sub><em>i</em></sub> is <tt>num_output_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_samples, ///< [in] The number of samples between starts of consecutive rows in the region of interest | |
| cudaStream_t stream, ///< [in] CUDA stream to launch kernels within. Default is stream<sub>0</sub>. | |
| bool debug_synchronous, ///< [in] Whether or not to synchronize the stream after every kernel launch to check for errors. May cause significant slowdown. Default is \p false. | |
| Int2Type<false> /*is_byte_sample*/) ///< [in] Marker type indicating whether or not SampleT is a 8b type | |
| { | |
| cudaError error = cudaSuccess; | |
| do | |
| { | |
| // Get PTX version | |
| int ptx_version = 0; | |
| if (CubDebug(error = PtxVersion(ptx_version))) break; | |
| // Get kernel dispatch configurations | |
| KernelConfig histogram_sweep_config; | |
| if (CubDebug(error = InitConfigs(ptx_version, histogram_sweep_config))) | |
| break; | |
| // Use the scale transform op for converting samples to privatized bins | |
| typedef ScaleTransform PrivatizedDecodeOpT; | |
| // Use the pass-thru transform op for converting privatized bins to output bins | |
| typedef PassThruTransform OutputDecodeOpT; | |
| PrivatizedDecodeOpT privatized_decode_op[NUM_ACTIVE_CHANNELS]; | |
| OutputDecodeOpT output_decode_op[NUM_ACTIVE_CHANNELS]; | |
| int max_levels = num_output_levels[0]; | |
| for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel) | |
| { | |
| int bins = num_output_levels[channel] - 1; | |
| LevelT scale = (upper_level[channel] - lower_level[channel]) / bins; | |
| privatized_decode_op[channel].Init(num_output_levels[channel], upper_level[channel], lower_level[channel], scale); | |
| if (num_output_levels[channel] > max_levels) | |
| max_levels = num_output_levels[channel]; | |
| } | |
| int max_num_output_bins = max_levels - 1; | |
| if (max_num_output_bins > MAX_PRIVATIZED_SMEM_BINS) | |
| { | |
| // Dispatch shared-privatized approach | |
| const int PRIVATIZED_SMEM_BINS = 0; | |
| if (CubDebug(error = PrivatizedDispatch( | |
| d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_output_histograms, | |
| num_output_levels, | |
| privatized_decode_op, | |
| num_output_levels, | |
| output_decode_op, | |
| max_num_output_bins, | |
| num_row_pixels, | |
| num_rows, | |
| row_stride_samples, | |
| DeviceHistogramInitKernel<NUM_ACTIVE_CHANNELS, CounterT, OffsetT>, | |
| DeviceHistogramSweepKernel<PtxHistogramSweepPolicy, PRIVATIZED_SMEM_BINS, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleIteratorT, CounterT, PrivatizedDecodeOpT, OutputDecodeOpT, OffsetT>, | |
| histogram_sweep_config, | |
| stream, | |
| debug_synchronous))) break; | |
| } | |
| else | |
| { | |
| // Dispatch shared-privatized approach | |
| const int PRIVATIZED_SMEM_BINS = MAX_PRIVATIZED_SMEM_BINS; | |
| if (CubDebug(error = PrivatizedDispatch( | |
| d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_output_histograms, | |
| num_output_levels, | |
| privatized_decode_op, | |
| num_output_levels, | |
| output_decode_op, | |
| max_num_output_bins, | |
| num_row_pixels, | |
| num_rows, | |
| row_stride_samples, | |
| DeviceHistogramInitKernel<NUM_ACTIVE_CHANNELS, CounterT, OffsetT>, | |
| DeviceHistogramSweepKernel<PtxHistogramSweepPolicy, PRIVATIZED_SMEM_BINS, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleIteratorT, CounterT, PrivatizedDecodeOpT, OutputDecodeOpT, OffsetT>, | |
| histogram_sweep_config, | |
| stream, | |
| debug_synchronous))) break; | |
| } | |
| } | |
| while (0); | |
| return error; | |
| } | |
| /** | |
| * Dispatch routine for HistogramEven, specialized for 8-bit sample types (computes 256-bin privatized histograms and then reduces to user-specified levels) | |
| */ | |
| CUB_RUNTIME_FUNCTION __forceinline__ | |
| static cudaError_t DispatchEven( | |
| void* d_temp_storage, ///< [in] %Device-accessible allocation of temporary storage. When NULL, the required allocation size is written to \p temp_storage_bytes and no work is done. | |
| size_t& temp_storage_bytes, ///< [in,out] Reference to size in bytes of \p d_temp_storage allocation | |
| SampleIteratorT d_samples, ///< [in] The pointer to the input sequence of sample items. 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_output_histograms[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_output_levels[i]</tt> - 1. | |
| int num_output_levels[NUM_ACTIVE_CHANNELS], ///< [in] The number of bin level boundaries for delineating histogram samples in each active channel. Implies that the number of bins for channel<sub><em>i</em></sub> is <tt>num_output_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_samples, ///< [in] The number of samples between starts of consecutive rows in the region of interest | |
| cudaStream_t stream, ///< [in] CUDA stream to launch kernels within. Default is stream<sub>0</sub>. | |
| bool debug_synchronous, ///< [in] Whether or not to synchronize the stream after every kernel launch to check for errors. May cause significant slowdown. Default is \p false. | |
| Int2Type<true> /*is_byte_sample*/) ///< [in] Marker type indicating whether or not SampleT is a 8b type | |
| { | |
| cudaError error = cudaSuccess; | |
| do | |
| { | |
| // Get PTX version | |
| int ptx_version = 0; | |
| if (CubDebug(error = PtxVersion(ptx_version))) break; | |
| // Get kernel dispatch configurations | |
| KernelConfig histogram_sweep_config; | |
| if (CubDebug(error = InitConfigs(ptx_version, histogram_sweep_config))) | |
| break; | |
| // Use the pass-thru transform op for converting samples to privatized bins | |
| typedef PassThruTransform PrivatizedDecodeOpT; | |
| // Use the scale transform op for converting privatized bins to output bins | |
| typedef ScaleTransform OutputDecodeOpT; | |
| int num_privatized_levels[NUM_ACTIVE_CHANNELS]; | |
| PrivatizedDecodeOpT privatized_decode_op[NUM_ACTIVE_CHANNELS]; | |
| OutputDecodeOpT output_decode_op[NUM_ACTIVE_CHANNELS]; | |
| int max_levels = num_output_levels[0]; | |
| for (int channel = 0; channel < NUM_ACTIVE_CHANNELS; ++channel) | |
| { | |
| num_privatized_levels[channel] = 257; | |
| int bins = num_output_levels[channel] - 1; | |
| LevelT scale = (upper_level[channel] - lower_level[channel]) / bins; | |
| output_decode_op[channel].Init(num_output_levels[channel], upper_level[channel], lower_level[channel], scale); | |
| if (num_output_levels[channel] > max_levels) | |
| max_levels = num_output_levels[channel]; | |
| } | |
| int max_num_output_bins = max_levels - 1; | |
| const int PRIVATIZED_SMEM_BINS = 256; | |
| if (CubDebug(error = PrivatizedDispatch( | |
| d_temp_storage, | |
| temp_storage_bytes, | |
| d_samples, | |
| d_output_histograms, | |
| num_privatized_levels, | |
| privatized_decode_op, | |
| num_output_levels, | |
| output_decode_op, | |
| max_num_output_bins, | |
| num_row_pixels, | |
| num_rows, | |
| row_stride_samples, | |
| DeviceHistogramInitKernel<NUM_ACTIVE_CHANNELS, CounterT, OffsetT>, | |
| DeviceHistogramSweepKernel<PtxHistogramSweepPolicy, PRIVATIZED_SMEM_BINS, NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleIteratorT, CounterT, PrivatizedDecodeOpT, OutputDecodeOpT, OffsetT>, | |
| histogram_sweep_config, | |
| stream, | |
| debug_synchronous))) break; | |
| } | |
| while (0); | |
| return error; | |
| } | |
| }; | |
| } // CUB namespace | |
| CUB_NS_POSTFIX // Optional outer namespace(s) | |