| /****************************************************************************** |
| * 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::AgentHistogram implements a stateful abstraction of CUDA thread blocks for participating in device-wide histogram . |
| */ |
| |
| #pragma once |
| |
| #include <iterator> |
| |
| #include "../util_type.cuh" |
| #include "../block/block_load.cuh" |
| #include "../config.cuh" |
| #include "../grid/grid_queue.cuh" |
| #include "../iterator/cache_modified_input_iterator.cuh" |
| |
| CUB_NAMESPACE_BEGIN |
| |
| |
| /****************************************************************************** |
| * Tuning policy |
| ******************************************************************************/ |
| |
| /** |
| * |
| */ |
| enum BlockHistogramMemoryPreference |
| { |
| GMEM, |
| SMEM, |
| BLEND |
| }; |
| |
| |
| /** |
| * Parameterizable tuning policy type for AgentHistogram |
| */ |
| template < |
| int _BLOCK_THREADS, ///< Threads per thread block |
| int _PIXELS_PER_THREAD, ///< Pixels per thread (per tile of input) |
| BlockLoadAlgorithm _LOAD_ALGORITHM, ///< The BlockLoad algorithm to use |
| CacheLoadModifier _LOAD_MODIFIER, ///< Cache load modifier for reading input elements |
| bool _RLE_COMPRESS, ///< Whether to perform localized RLE to compress samples before histogramming |
| BlockHistogramMemoryPreference _MEM_PREFERENCE, ///< Whether to prefer privatized shared-memory bins (versus privatized global-memory bins) |
| bool _WORK_STEALING> ///< Whether to dequeue tiles from a global work queue |
| struct AgentHistogramPolicy |
| { |
| enum |
| { |
| BLOCK_THREADS = _BLOCK_THREADS, ///< Threads per thread block |
| PIXELS_PER_THREAD = _PIXELS_PER_THREAD, ///< Pixels per thread (per tile of input) |
| IS_RLE_COMPRESS = _RLE_COMPRESS, ///< Whether to perform localized RLE to compress samples before histogramming |
| MEM_PREFERENCE = _MEM_PREFERENCE, ///< Whether to prefer privatized shared-memory bins (versus privatized global-memory bins) |
| IS_WORK_STEALING = _WORK_STEALING, ///< Whether to dequeue tiles from a global work queue |
| }; |
| |
| static const BlockLoadAlgorithm LOAD_ALGORITHM = _LOAD_ALGORITHM; ///< The BlockLoad algorithm to use |
| static const CacheLoadModifier LOAD_MODIFIER = _LOAD_MODIFIER; ///< Cache load modifier for reading input elements |
| }; |
| |
| |
| /****************************************************************************** |
| * Thread block abstractions |
| ******************************************************************************/ |
| |
| /** |
| * \brief AgentHistogram implements a stateful abstraction of CUDA thread blocks for participating in device-wide histogram . |
| */ |
| template < |
| typename AgentHistogramPolicyT, ///< Parameterized AgentHistogramPolicy tuning policy type |
| int PRIVATIZED_SMEM_BINS, ///< Number of privatized shared-memory histogram bins of any channel. Zero indicates privatized counters to be maintained in device-accessible memory. |
| int NUM_CHANNELS, ///< Number of channels interleaved in the input data. Supports up to four channels. |
| int NUM_ACTIVE_CHANNELS, ///< Number of channels actively being histogrammed |
| typename SampleIteratorT, ///< Random-access input iterator type for reading samples |
| 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 |
| int LEGACY_PTX_ARCH = 0> ///< PTX compute capability (unused) |
| struct AgentHistogram |
| { |
| //--------------------------------------------------------------------- |
| // Types and constants |
| //--------------------------------------------------------------------- |
| |
| /// The sample type of the input iterator |
| using SampleT = cub::detail::value_t<SampleIteratorT>; |
| |
| /// The pixel type of SampleT |
| using PixelT = typename CubVector<SampleT, NUM_CHANNELS>::Type; |
| |
| /// The quad type of SampleT |
| using QuadT = typename CubVector<SampleT, 4>::Type; |
| |
| /// Constants |
| enum |
| { |
| BLOCK_THREADS = AgentHistogramPolicyT::BLOCK_THREADS, |
| |
| PIXELS_PER_THREAD = AgentHistogramPolicyT::PIXELS_PER_THREAD, |
| SAMPLES_PER_THREAD = PIXELS_PER_THREAD * NUM_CHANNELS, |
| QUADS_PER_THREAD = SAMPLES_PER_THREAD / 4, |
| |
| TILE_PIXELS = PIXELS_PER_THREAD * BLOCK_THREADS, |
| TILE_SAMPLES = SAMPLES_PER_THREAD * BLOCK_THREADS, |
| |
| IS_RLE_COMPRESS = AgentHistogramPolicyT::IS_RLE_COMPRESS, |
| |
| MEM_PREFERENCE = (PRIVATIZED_SMEM_BINS > 0) ? |
| AgentHistogramPolicyT::MEM_PREFERENCE : |
| GMEM, |
| |
| IS_WORK_STEALING = AgentHistogramPolicyT::IS_WORK_STEALING, |
| }; |
| |
| /// Cache load modifier for reading input elements |
| static const CacheLoadModifier LOAD_MODIFIER = AgentHistogramPolicyT::LOAD_MODIFIER; |
| |
|
|
| /// Input iterator wrapper type (for applying cache modifier) |
| // Wrap the native input pointer with CacheModifiedInputIterator |
| // or directly use the supplied input iterator type |
| using WrappedSampleIteratorT = cub::detail::conditional_t< |
| std::is_pointer<SampleIteratorT>::value, |
| CacheModifiedInputIterator<LOAD_MODIFIER, SampleT, OffsetT>, |
| SampleIteratorT>; |
| |
| /// Pixel input iterator type (for applying cache modifier) |
| typedef CacheModifiedInputIterator<LOAD_MODIFIER, PixelT, OffsetT> |
| WrappedPixelIteratorT; |
| |
| /// Qaud input iterator type (for applying cache modifier) |
| typedef CacheModifiedInputIterator<LOAD_MODIFIER, QuadT, OffsetT> |
| WrappedQuadIteratorT; |
| |
| /// Parameterized BlockLoad type for samples |
| typedef BlockLoad< |
| SampleT, |
| BLOCK_THREADS, |
| SAMPLES_PER_THREAD, |
| AgentHistogramPolicyT::LOAD_ALGORITHM> |
| BlockLoadSampleT; |
| |
| /// Parameterized BlockLoad type for pixels |
| typedef BlockLoad< |
| PixelT, |
| BLOCK_THREADS, |
| PIXELS_PER_THREAD, |
| AgentHistogramPolicyT::LOAD_ALGORITHM> |
| BlockLoadPixelT; |
| |
| /// Parameterized BlockLoad type for quads |
| typedef BlockLoad< |
| QuadT, |
| BLOCK_THREADS, |
| QUADS_PER_THREAD, |
| AgentHistogramPolicyT::LOAD_ALGORITHM> |
| BlockLoadQuadT; |
| |
| /// Shared memory type required by this thread block |
| struct _TempStorage |
| { |
| CounterT histograms[NUM_ACTIVE_CHANNELS][PRIVATIZED_SMEM_BINS + 1]; // Smem needed for block-privatized smem histogram (with 1 word of padding) |
| |
| int tile_idx; |
| |
| // Aliasable storage layout |
| union Aliasable |
| { |
| typename BlockLoadSampleT::TempStorage sample_load; // Smem needed for loading a tile of samples |
| typename BlockLoadPixelT::TempStorage pixel_load; // Smem needed for loading a tile of pixels |
| typename BlockLoadQuadT::TempStorage quad_load; // Smem needed for loading a tile of quads |
| |
| } aliasable; |
| }; |
| |
|
|
| /// Temporary storage type (unionable) |
| struct TempStorage : Uninitialized<_TempStorage> {}; |
| |
|
|
| //--------------------------------------------------------------------- |
| // Per-thread fields |
| //--------------------------------------------------------------------- |
| |
| /// Reference to temp_storage |
| _TempStorage &temp_storage; |
| |
| /// Sample input iterator (with cache modifier applied, if possible) |
| WrappedSampleIteratorT d_wrapped_samples; |
| |
| /// Native pointer for input samples (possibly NULL if unavailable) |
| SampleT* d_native_samples; |
| |
| /// The number of output bins for each channel |
| int (&num_output_bins)[NUM_ACTIVE_CHANNELS]; |
| |
| /// The number of privatized bins for each channel |
| int (&num_privatized_bins)[NUM_ACTIVE_CHANNELS]; |
| |
| /// Reference to gmem privatized histograms for each channel |
| CounterT* d_privatized_histograms[NUM_ACTIVE_CHANNELS]; |
| |
| /// Reference to final output histograms (gmem) |
| CounterT* (&d_output_histograms)[NUM_ACTIVE_CHANNELS]; |
| |
| /// The transform operator for determining output bin-ids from privatized counter indices, one for each channel |
| OutputDecodeOpT (&output_decode_op)[NUM_ACTIVE_CHANNELS]; |
| |
| /// The transform operator for determining privatized counter indices from samples, one for each channel |
| PrivatizedDecodeOpT (&privatized_decode_op)[NUM_ACTIVE_CHANNELS]; |
| |
| /// Whether to prefer privatized smem counters vs privatized global counters |
| bool prefer_smem; |
| |
|
|
| //--------------------------------------------------------------------- |
| // Initialize privatized bin counters |
| //--------------------------------------------------------------------- |
| |
| // Initialize privatized bin counters |
| __device__ __forceinline__ void InitBinCounters(CounterT* privatized_histograms[NUM_ACTIVE_CHANNELS]) |
| { |
| // Initialize histogram bin counts to zeros |
| #pragma unroll |
| for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL) |
| { |
| for (int privatized_bin = threadIdx.x; privatized_bin < num_privatized_bins[CHANNEL]; privatized_bin += BLOCK_THREADS) |
| { |
| privatized_histograms[CHANNEL][privatized_bin] = 0; |
| } |
| } |
| |
| // Barrier to make sure all threads are done updating counters |
| CTA_SYNC(); |
| } |
| |
|
|
| // Initialize privatized bin counters. Specialized for privatized shared-memory counters |
| __device__ __forceinline__ void InitSmemBinCounters() |
| { |
| CounterT* privatized_histograms[NUM_ACTIVE_CHANNELS]; |
| |
| for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL) |
| privatized_histograms[CHANNEL] = temp_storage.histograms[CHANNEL]; |
| |
| InitBinCounters(privatized_histograms); |
| } |
| |
|
|
| // Initialize privatized bin counters. Specialized for privatized global-memory counters |
| __device__ __forceinline__ void InitGmemBinCounters() |
| { |
| InitBinCounters(d_privatized_histograms); |
| } |
| |
|
|
| //--------------------------------------------------------------------- |
| // Update final output histograms |
| //--------------------------------------------------------------------- |
| |
| // Update final output histograms from privatized histograms |
| __device__ __forceinline__ void StoreOutput(CounterT* privatized_histograms[NUM_ACTIVE_CHANNELS]) |
| { |
| // Barrier to make sure all threads are done updating counters |
| CTA_SYNC(); |
| |
| // Apply privatized bin counts to output bin counts |
| #pragma unroll |
| for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL) |
| { |
| int channel_bins = num_privatized_bins[CHANNEL]; |
| for (int privatized_bin = threadIdx.x; |
| privatized_bin < channel_bins; |
| privatized_bin += BLOCK_THREADS) |
| { |
| int output_bin = -1; |
| CounterT count = privatized_histograms[CHANNEL][privatized_bin]; |
| bool is_valid = count > 0; |
| |
| output_decode_op[CHANNEL].template BinSelect<LOAD_MODIFIER>((SampleT) privatized_bin, output_bin, is_valid); |
| |
| if (output_bin >= 0) |
| { |
| atomicAdd(&d_output_histograms[CHANNEL][output_bin], count); |
| } |
| |
| } |
| } |
| } |
| |
|
|
| // Update final output histograms from privatized histograms. Specialized for privatized shared-memory counters |
| __device__ __forceinline__ void StoreSmemOutput() |
| { |
| CounterT* privatized_histograms[NUM_ACTIVE_CHANNELS]; |
| for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL) |
| privatized_histograms[CHANNEL] = temp_storage.histograms[CHANNEL]; |
| |
| StoreOutput(privatized_histograms); |
| } |
| |
|
|
| // Update final output histograms from privatized histograms. Specialized for privatized global-memory counters |
| __device__ __forceinline__ void StoreGmemOutput() |
| { |
| StoreOutput(d_privatized_histograms); |
| } |
| |
|
|
| //--------------------------------------------------------------------- |
| // Tile accumulation |
| //--------------------------------------------------------------------- |
| |
| // Accumulate pixels. Specialized for RLE compression. |
| __device__ __forceinline__ void AccumulatePixels( |
| SampleT samples[PIXELS_PER_THREAD][NUM_CHANNELS], |
| bool is_valid[PIXELS_PER_THREAD], |
| CounterT* privatized_histograms[NUM_ACTIVE_CHANNELS], |
| Int2Type<true> is_rle_compress) |
| { |
| #pragma unroll |
| for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL) |
| { |
| // Bin pixels |
| int bins[PIXELS_PER_THREAD]; |
| |
| #pragma unroll |
| for (int PIXEL = 0; PIXEL < PIXELS_PER_THREAD; ++PIXEL) |
| { |
| bins[PIXEL] = -1; |
| privatized_decode_op[CHANNEL].template BinSelect<LOAD_MODIFIER>(samples[PIXEL][CHANNEL], bins[PIXEL], is_valid[PIXEL]); |
| } |
| |
| CounterT accumulator = 1; |
| |
| #pragma unroll |
| for (int PIXEL = 0; PIXEL < PIXELS_PER_THREAD - 1; ++PIXEL) |
| { |
| if (bins[PIXEL] != bins[PIXEL + 1]) |
| { |
| if (bins[PIXEL] >= 0) |
| atomicAdd(privatized_histograms[CHANNEL] + bins[PIXEL], accumulator); |
| |
| accumulator = 0; |
| } |
| accumulator++; |
| } |
| |
| // Last pixel |
| if (bins[PIXELS_PER_THREAD - 1] >= 0) |
| atomicAdd(privatized_histograms[CHANNEL] + bins[PIXELS_PER_THREAD - 1], accumulator); |
| } |
| } |
| |
|
|
| // Accumulate pixels. Specialized for individual accumulation of each pixel. |
| __device__ __forceinline__ void AccumulatePixels( |
| SampleT samples[PIXELS_PER_THREAD][NUM_CHANNELS], |
| bool is_valid[PIXELS_PER_THREAD], |
| CounterT* privatized_histograms[NUM_ACTIVE_CHANNELS], |
| Int2Type<false> is_rle_compress) |
| { |
| #pragma unroll |
| for (int PIXEL = 0; PIXEL < PIXELS_PER_THREAD; ++PIXEL) |
| { |
| #pragma unroll |
| for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL) |
| { |
| int bin = -1; |
| privatized_decode_op[CHANNEL].template BinSelect<LOAD_MODIFIER>(samples[PIXEL][CHANNEL], bin, is_valid[PIXEL]); |
| if (bin >= 0) |
| atomicAdd(privatized_histograms[CHANNEL] + bin, 1); |
| } |
| } |
| } |
| |
|
|
| /** |
| * Accumulate pixel, specialized for smem privatized histogram |
| */ |
| __device__ __forceinline__ void AccumulateSmemPixels( |
| SampleT samples[PIXELS_PER_THREAD][NUM_CHANNELS], |
| bool is_valid[PIXELS_PER_THREAD]) |
| { |
| CounterT* privatized_histograms[NUM_ACTIVE_CHANNELS]; |
| |
| for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL) |
| privatized_histograms[CHANNEL] = temp_storage.histograms[CHANNEL]; |
| |
| AccumulatePixels(samples, is_valid, privatized_histograms, Int2Type<IS_RLE_COMPRESS>()); |
| } |
| |
|
|
| /** |
| * Accumulate pixel, specialized for gmem privatized histogram |
| */ |
| __device__ __forceinline__ void AccumulateGmemPixels( |
| SampleT samples[PIXELS_PER_THREAD][NUM_CHANNELS], |
| bool is_valid[PIXELS_PER_THREAD]) |
| { |
| AccumulatePixels(samples, is_valid, d_privatized_histograms, Int2Type<IS_RLE_COMPRESS>()); |
| } |
| |
|
|
|
|
| //--------------------------------------------------------------------- |
| // Tile loading |
| //--------------------------------------------------------------------- |
| |
| // Load full, aligned tile using pixel iterator (multi-channel) |
| template <int _NUM_ACTIVE_CHANNELS> |
| __device__ __forceinline__ void LoadFullAlignedTile( |
| OffsetT block_offset, |
| int valid_samples, |
| SampleT (&samples)[PIXELS_PER_THREAD][NUM_CHANNELS], |
| Int2Type<_NUM_ACTIVE_CHANNELS> num_active_channels) |
| { |
| typedef PixelT AliasedPixels[PIXELS_PER_THREAD]; |
| |
| WrappedPixelIteratorT d_wrapped_pixels((PixelT*) (d_native_samples + block_offset)); |
| |
| // Load using a wrapped pixel iterator |
| BlockLoadPixelT(temp_storage.aliasable.pixel_load).Load( |
| d_wrapped_pixels, |
| reinterpret_cast<AliasedPixels&>(samples)); |
| } |
| |
| // Load full, aligned tile using quad iterator (single-channel) |
| __device__ __forceinline__ void LoadFullAlignedTile( |
| OffsetT block_offset, |
| int valid_samples, |
| SampleT (&samples)[PIXELS_PER_THREAD][NUM_CHANNELS], |
| Int2Type<1> num_active_channels) |
| { |
| typedef QuadT AliasedQuads[QUADS_PER_THREAD]; |
| |
| WrappedQuadIteratorT d_wrapped_quads((QuadT*) (d_native_samples + block_offset)); |
| |
| // Load using a wrapped quad iterator |
| BlockLoadQuadT(temp_storage.aliasable.quad_load).Load( |
| d_wrapped_quads, |
| reinterpret_cast<AliasedQuads&>(samples)); |
| } |
| |
| // Load full, aligned tile |
| __device__ __forceinline__ void LoadTile( |
| OffsetT block_offset, |
| int valid_samples, |
| SampleT (&samples)[PIXELS_PER_THREAD][NUM_CHANNELS], |
| Int2Type<true> is_full_tile, |
| Int2Type<true> is_aligned) |
| { |
| LoadFullAlignedTile(block_offset, valid_samples, samples, Int2Type<NUM_ACTIVE_CHANNELS>()); |
| } |
| |
| // Load full, mis-aligned tile using sample iterator |
| __device__ __forceinline__ void LoadTile( |
| OffsetT block_offset, |
| int valid_samples, |
| SampleT (&samples)[PIXELS_PER_THREAD][NUM_CHANNELS], |
| Int2Type<true> is_full_tile, |
| Int2Type<false> is_aligned) |
| { |
| typedef SampleT AliasedSamples[SAMPLES_PER_THREAD]; |
| |
| // Load using sample iterator |
| BlockLoadSampleT(temp_storage.aliasable.sample_load).Load( |
| d_wrapped_samples + block_offset, |
| reinterpret_cast<AliasedSamples&>(samples)); |
| } |
| |
| // Load partially-full, aligned tile using the pixel iterator |
| __device__ __forceinline__ void LoadTile( |
| OffsetT block_offset, |
| int valid_samples, |
| SampleT (&samples)[PIXELS_PER_THREAD][NUM_CHANNELS], |
| Int2Type<false> is_full_tile, |
| Int2Type<true> is_aligned) |
| { |
| typedef PixelT AliasedPixels[PIXELS_PER_THREAD]; |
| |
| WrappedPixelIteratorT d_wrapped_pixels((PixelT*) (d_native_samples + block_offset)); |
| |
| int valid_pixels = valid_samples / NUM_CHANNELS; |
| |
| // Load using a wrapped pixel iterator |
| BlockLoadPixelT(temp_storage.aliasable.pixel_load).Load( |
| d_wrapped_pixels, |
| reinterpret_cast<AliasedPixels&>(samples), |
| valid_pixels); |
| } |
| |
| // Load partially-full, mis-aligned tile using sample iterator |
| __device__ __forceinline__ void LoadTile( |
| OffsetT block_offset, |
| int valid_samples, |
| SampleT (&samples)[PIXELS_PER_THREAD][NUM_CHANNELS], |
| Int2Type<false> is_full_tile, |
| Int2Type<false> is_aligned) |
| { |
| typedef SampleT AliasedSamples[SAMPLES_PER_THREAD]; |
| |
| BlockLoadSampleT(temp_storage.aliasable.sample_load).Load( |
| d_wrapped_samples + block_offset, |
| reinterpret_cast<AliasedSamples&>(samples), |
| valid_samples); |
| } |
| |
|
|
| //--------------------------------------------------------------------- |
| // Tile processing |
| //--------------------------------------------------------------------- |
| |
| // Consume a tile of data samples |
| template < |
| bool IS_ALIGNED, // Whether the tile offset is aligned (quad-aligned for single-channel, pixel-aligned for multi-channel) |
| bool IS_FULL_TILE> // Whether the tile is full |
| __device__ __forceinline__ void ConsumeTile(OffsetT block_offset, int valid_samples) |
| { |
| SampleT samples[PIXELS_PER_THREAD][NUM_CHANNELS]; |
| bool is_valid[PIXELS_PER_THREAD]; |
| |
| // Load tile |
| LoadTile( |
| block_offset, |
| valid_samples, |
| samples, |
| Int2Type<IS_FULL_TILE>(), |
| Int2Type<IS_ALIGNED>()); |
| |
| // Set valid flags |
| #pragma unroll |
| for (int PIXEL = 0; PIXEL < PIXELS_PER_THREAD; ++PIXEL) |
| is_valid[PIXEL] = IS_FULL_TILE || (((threadIdx.x * PIXELS_PER_THREAD + PIXEL) * NUM_CHANNELS) < valid_samples); |
| |
| // Accumulate samples |
| if (prefer_smem) |
| AccumulateSmemPixels(samples, is_valid); |
| else |
| AccumulateGmemPixels(samples, is_valid); |
| } |
| |
|
|
| // Consume row tiles. Specialized for work-stealing from queue |
| template <bool IS_ALIGNED> |
| __device__ __forceinline__ void ConsumeTiles( |
| 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, |
| Int2Type<true> is_work_stealing) |
| { |
| |
| int num_tiles = num_rows * tiles_per_row; |
| int tile_idx = (blockIdx.y * gridDim.x) + blockIdx.x; |
| OffsetT num_even_share_tiles = gridDim.x * gridDim.y; |
| |
| while (tile_idx < num_tiles) |
| { |
| int row = tile_idx / tiles_per_row; |
| int col = tile_idx - (row * tiles_per_row); |
| OffsetT row_offset = row * row_stride_samples; |
| OffsetT col_offset = (col * TILE_SAMPLES); |
| OffsetT tile_offset = row_offset + col_offset; |
| |
| if (col == tiles_per_row - 1) |
| { |
| // Consume a partially-full tile at the end of the row |
| OffsetT num_remaining = (num_row_pixels * NUM_CHANNELS) - col_offset; |
| ConsumeTile<IS_ALIGNED, false>(tile_offset, num_remaining); |
| } |
| else |
| { |
| // Consume full tile |
| ConsumeTile<IS_ALIGNED, true>(tile_offset, TILE_SAMPLES); |
| } |
| |
| CTA_SYNC(); |
| |
| // Get next tile |
| if (threadIdx.x == 0) |
| temp_storage.tile_idx = tile_queue.Drain(1) + num_even_share_tiles; |
| |
| CTA_SYNC(); |
| |
| tile_idx = temp_storage.tile_idx; |
| } |
| } |
| |
|
|
| // Consume row tiles. Specialized for even-share (striped across thread blocks) |
| template <bool IS_ALIGNED> |
| __device__ __forceinline__ void ConsumeTiles( |
| 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, |
| Int2Type<false> is_work_stealing) |
| { |
| for (int row = blockIdx.y; row < num_rows; row += gridDim.y) |
| { |
| OffsetT row_begin = row * row_stride_samples; |
| OffsetT row_end = row_begin + (num_row_pixels * NUM_CHANNELS); |
| OffsetT tile_offset = row_begin + (blockIdx.x * TILE_SAMPLES); |
| |
| while (tile_offset < row_end) |
| { |
| OffsetT num_remaining = row_end - tile_offset; |
| |
| if (num_remaining < TILE_SAMPLES) |
| { |
| // Consume partial tile |
| ConsumeTile<IS_ALIGNED, false>(tile_offset, num_remaining); |
| break; |
| } |
| |
| // Consume full tile |
| ConsumeTile<IS_ALIGNED, true>(tile_offset, TILE_SAMPLES); |
| tile_offset += gridDim.x * TILE_SAMPLES; |
| } |
| } |
| } |
| |
|
|
| //--------------------------------------------------------------------- |
| // Parameter extraction |
| //--------------------------------------------------------------------- |
| |
| // Return a native pixel pointer (specialized for CacheModifiedInputIterator types) |
| template < |
| CacheLoadModifier _MODIFIER, |
| typename _ValueT, |
| typename _OffsetT> |
| __device__ __forceinline__ SampleT* NativePointer(CacheModifiedInputIterator<_MODIFIER, _ValueT, _OffsetT> itr) |
| { |
| return itr.ptr; |
| } |
| |
| // Return a native pixel pointer (specialized for other types) |
| template <typename IteratorT> |
| __device__ __forceinline__ SampleT* NativePointer(IteratorT itr) |
| { |
| return NULL; |
| } |
| |
|
|
|
|
| //--------------------------------------------------------------------- |
| // Interface |
| //--------------------------------------------------------------------- |
| |
|
|
| /** |
| * Constructor |
| */ |
| __device__ __forceinline__ AgentHistogram( |
| TempStorage &temp_storage, ///< Reference to temp_storage |
| SampleIteratorT d_samples, ///< Input data to reduce |
| int (&num_output_bins)[NUM_ACTIVE_CHANNELS], ///< The number bins per final output histogram |
| int (&num_privatized_bins)[NUM_ACTIVE_CHANNELS], ///< The number bins per privatized histogram |
| CounterT* (&d_output_histograms)[NUM_ACTIVE_CHANNELS], ///< Reference to final output histograms |
| CounterT* (&d_privatized_histograms)[NUM_ACTIVE_CHANNELS], ///< Reference to privatized histograms |
| OutputDecodeOpT (&output_decode_op)[NUM_ACTIVE_CHANNELS], ///< The transform operator for determining output bin-ids from privatized counter indices, one for each channel |
| PrivatizedDecodeOpT (&privatized_decode_op)[NUM_ACTIVE_CHANNELS]) ///< The transform operator for determining privatized counter indices from samples, one for each channel |
| : |
| temp_storage(temp_storage.Alias()), |
| d_wrapped_samples(d_samples), |
| d_native_samples(NativePointer(d_wrapped_samples)), |
| num_output_bins(num_output_bins), |
| num_privatized_bins(num_privatized_bins), |
| d_output_histograms(d_output_histograms), |
| output_decode_op(output_decode_op), |
| privatized_decode_op(privatized_decode_op), |
| prefer_smem((MEM_PREFERENCE == SMEM) ? |
| true : // prefer smem privatized histograms |
| (MEM_PREFERENCE == GMEM) ? |
| false : // prefer gmem privatized histograms |
| blockIdx.x & 1) // prefer blended privatized histograms |
| { |
| int blockId = (blockIdx.y * gridDim.x) + blockIdx.x; |
| |
| // Initialize the locations of this block's privatized histograms |
| for (int CHANNEL = 0; CHANNEL < NUM_ACTIVE_CHANNELS; ++CHANNEL) |
| this->d_privatized_histograms[CHANNEL] = d_privatized_histograms[CHANNEL] + (blockId * num_privatized_bins[CHANNEL]); |
| } |
| |
|
|
| /** |
| * Consume image |
| */ |
| __device__ __forceinline__ void ConsumeTiles( |
| 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) ///< Queue descriptor for assigning tiles of work to thread blocks |
| { |
| // Check whether all row starting offsets are quad-aligned (in single-channel) or pixel-aligned (in multi-channel) |
| int quad_mask = AlignBytes<QuadT>::ALIGN_BYTES - 1; |
| int pixel_mask = AlignBytes<PixelT>::ALIGN_BYTES - 1; |
| size_t row_bytes = sizeof(SampleT) * row_stride_samples; |
| |
| bool quad_aligned_rows = (NUM_CHANNELS == 1) && (SAMPLES_PER_THREAD % 4 == 0) && // Single channel |
| ((size_t(d_native_samples) & quad_mask) == 0) && // ptr is quad-aligned |
| ((num_rows == 1) || ((row_bytes & quad_mask) == 0)); // number of row-samples is a multiple of the alignment of the quad |
| |
| bool pixel_aligned_rows = (NUM_CHANNELS > 1) && // Multi channel |
| ((size_t(d_native_samples) & pixel_mask) == 0) && // ptr is pixel-aligned |
| ((row_bytes & pixel_mask) == 0); // number of row-samples is a multiple of the alignment of the pixel |
| |
| // Whether rows are aligned and can be vectorized |
| if ((d_native_samples != NULL) && (quad_aligned_rows || pixel_aligned_rows)) |
| ConsumeTiles<true>(num_row_pixels, num_rows, row_stride_samples, tiles_per_row, tile_queue, Int2Type<IS_WORK_STEALING>()); |
| else |
| ConsumeTiles<false>(num_row_pixels, num_rows, row_stride_samples, tiles_per_row, tile_queue, Int2Type<IS_WORK_STEALING>()); |
| } |
| |
|
|
| /** |
| * Initialize privatized bin counters. Specialized for privatized shared-memory counters |
| */ |
| __device__ __forceinline__ void InitBinCounters() |
| { |
| if (prefer_smem) |
| InitSmemBinCounters(); |
| else |
| InitGmemBinCounters(); |
| } |
| |
|
|
| /** |
| * Store privatized histogram to device-accessible memory. Specialized for privatized shared-memory counters |
| */ |
| __device__ __forceinline__ void StoreOutput() |
| { |
| if (prefer_smem) |
| StoreSmemOutput(); |
| else |
| StoreGmemOutput(); |
| } |
| |
|
|
| }; |
|
|
| CUB_NAMESPACE_END |
|
|
