/****************************************************************************** * 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 #include #include #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 /// 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 num_output_bins_wrapper, ///< Number of output histogram bins per channel ArrayWrapper d_output_histograms_wrapper, ///< Histogram counter data having logical dimensions CounterT[NUM_ACTIVE_CHANNELS][num_bins.array[CHANNEL]] GridQueue 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 num_output_bins_wrapper, ///< The number bins per final output histogram ArrayWrapper num_privatized_bins_wrapper, ///< The number bins per privatized histogram ArrayWrapper d_output_histograms_wrapper, ///< Reference to final output histograms ArrayWrapper d_privatized_histograms_wrapper, ///< Reference to privatized histograms ArrayWrapper output_decode_op_wrapper, ///< The transform operator for determining output bin-ids from privatized counter indices, one for each channel ArrayWrapper 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 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::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 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 __host__ __device__ __forceinline__ void BinSelect(_SampleT sample, int &bin, bool valid) { /// Level iterator wrapper type typedef typename If::VALUE, CacheModifiedInputIterator, // 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 __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 __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 __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 __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 __host__ __device__ __forceinline__ void BinSelect(_SampleT sample, int &bin, bool valid) { if (valid) bin = (int) sample; } }; //--------------------------------------------------------------------- // Tuning policies //--------------------------------------------------------------------- template 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 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(); #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(); } else if (ptx_version >= 350) { result = histogram_sweep_config.template Init(); } else if (ptx_version >= 300) { result = histogram_sweep_config.template Init(); } else if (ptx_version >= 200) { result = histogram_sweep_config.template Init(); } else { result = histogram_sweep_config.template Init(); } #endif } return result; } /** * Kernel kernel dispatch configuration */ struct KernelConfig { int block_threads; int pixels_per_thread; template 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_i, the allocation length of d_histograms[i] should be num_output_levels[i] - 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_i is num_output_levels[i] - 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_i is num_output_levels[i] - 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₀. 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::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 tile_queue(allocations[NUM_ALLOCATIONS - 1]); // Setup array wrapper for histogram channel output (because we can't pass static arrays as kernel parameters) ArrayWrapper 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 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 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 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 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 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_i, the allocation length of d_histograms[i] should be num_output_levels[i] - 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_i is num_output_levels[i] - 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₀. 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 /*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 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, DeviceHistogramSweepKernel, 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, DeviceHistogramSweepKernel, 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_i, the allocation length of d_histograms[i] should be num_output_levels[i] - 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_i is num_output_levels[i] - 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₀. 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 /*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 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, DeviceHistogramSweepKernel, 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_i, the allocation length of d_histograms[i] should be num_output_levels[i] - 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_i is num_output_levels[i] - 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₀. 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 /*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, DeviceHistogramSweepKernel, 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, DeviceHistogramSweepKernel, 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_i, the allocation length of d_histograms[i] should be num_output_levels[i] - 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_i is num_output_levels[i] - 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₀. 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 /*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, DeviceHistogramSweepKernel, histogram_sweep_config, stream, debug_synchronous))) break; } while (0); return error; } }; } // CUB namespace CUB_NS_POSTFIX // Optional outer namespace(s)