1
/******************************************************************************
2
* Copyright (c) 2011, Duane Merrill. All rights reserved.
3
* Copyright (c) 2011-2013, NVIDIA CORPORATION. All rights reserved.
5
* Redistribution and use in source and binary forms, with or without
6
* modification, are permitted provided that the following conditions are met:
7
* * Redistributions of source code must retain the above copyright
8
* notice, this list of conditions and the following disclaimer.
9
* * Redistributions in binary form must reproduce the above copyright
10
* notice, this list of conditions and the following disclaimer in the
11
* documentation and/or other materials provided with the distribution.
12
* * Neither the name of the NVIDIA CORPORATION nor the
13
* names of its contributors may be used to endorse or promote products
14
* derived from this software without specific prior written permission.
16
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
17
* ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
18
* WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
19
* DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
20
* DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
21
* (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
22
* LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
23
* ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
24
* (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
25
* SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
27
******************************************************************************/
31
* cub::BlockRakingLayout provides a conflict-free shared memory layout abstraction for warp-raking across thread block data.
37
#include "../util_macro.cuh"
38
#include "../util_arch.cuh"
39
#include "../util_namespace.cuh"
41
/// Optional outer namespace(s)
48
* \brief BlockRakingLayout provides a conflict-free shared memory layout abstraction for raking across thread block data. ![](raking.png)
49
* \ingroup BlockModule
52
* This type facilitates a shared memory usage pattern where a block of CUDA
53
* threads places elements into shared memory and then reduces the active
54
* parallelism to one "raking" warp of threads for serially aggregating consecutive
55
* sequences of shared items. Padding is inserted to eliminate bank conflicts
56
* (for most data types).
58
* \tparam T The data type to be exchanged.
59
* \tparam BLOCK_THREADS The thread block size in threads.
60
* \tparam BLOCK_STRIPS When strip-mining, the number of threadblock-strips per tile
66
struct BlockRakingLayout
68
//---------------------------------------------------------------------
69
// Constants and typedefs
70
//---------------------------------------------------------------------
74
/// The total number of elements that need to be cooperatively reduced
76
BLOCK_THREADS * BLOCK_STRIPS,
78
/// Maximum number of warp-synchronous raking threads
80
CUB_MIN(BLOCK_THREADS, PtxArchProps::WARP_THREADS),
82
/// Number of raking elements per warp-synchronous raking thread (rounded up)
84
(SHARED_ELEMENTS + MAX_RAKING_THREADS - 1) / MAX_RAKING_THREADS,
86
/// Never use a raking thread that will have no valid data (e.g., when BLOCK_THREADS is 62 and SEGMENT_LENGTH is 2, we should only use 31 raking threads)
88
(SHARED_ELEMENTS + SEGMENT_LENGTH - 1) / SEGMENT_LENGTH,
90
/// Pad each segment length with one element if it evenly divides the number of banks
92
(PtxArchProps::SMEM_BANKS % SEGMENT_LENGTH == 0) ? 1 : 0,
94
/// Total number of elements in the raking grid
96
RAKING_THREADS * (SEGMENT_LENGTH + SEGMENT_PADDING),
98
/// Whether or not we need bounds checking during raking (the number of reduction elements is not a multiple of the warp size)
100
(SHARED_ELEMENTS % RAKING_THREADS == 0),
105
* \brief Shared memory storage type
107
typedef T TempStorage[BlockRakingLayout::GRID_ELEMENTS];
111
* \brief Returns the location for the calling thread to place data into the grid
113
static __device__ __forceinline__ T* PlacementPtr(
114
TempStorage &temp_storage,
118
// Offset for partial
119
unsigned int offset = (block_strip * BLOCK_THREADS) + linear_tid;
121
// Add in one padding element for every segment
122
if (SEGMENT_PADDING > 0)
124
offset += offset / SEGMENT_LENGTH;
127
// Incorporating a block of padding partials every shared memory segment
128
return temp_storage + offset;
133
* \brief Returns the location for the calling thread to begin sequential raking
135
static __device__ __forceinline__ T* RakingPtr(
136
TempStorage &temp_storage,
139
return temp_storage + (linear_tid * (SEGMENT_LENGTH + SEGMENT_PADDING));
144
CUB_NS_POSTFIX // Optional outer namespace(s)