CUB
|
The WarpReduce class provides collective methods for computing a parallel reduction of items partitioned across a CUDA thread warp.
T | The reduction input/output element type |
LOGICAL_WARP_THREADS | [optional] The number of threads per "logical" warp (may be less than the number of hardware warp threads). Default is the warp size of the targeted CUDA compute-capability (e.g., 32 threads for SM20). |
PTX_ARCH | [optional] The PTX compute capability for which to to specialize this collective, formatted as per the CUDA_ARCH macro (e.g., 350 for sm_35). Useful for determining the collective's storage requirements for a given device from the host. (Default: the value of CUDA_ARCH during the current compiler pass) |
LOGICAL_WARP_THREADS
SHFL
instructions)LOGICAL_WARP_THREADS
thread_data
across the block of threads is {0, 1, 2, 3, ..., 127}
. The corresponding output aggregate
in threads 0, 32, 64, and 96 will 496
, 1520
, 2544
, and 3568
, respectively (and is undefined in other threads).thread_data
across the warp of threads is {0, 1, 2, 3, ..., 31}
. The corresponding output aggregate
in thread0 will be 496
(and is undefined in other threads). Definition at line 141 of file warp_reduce.cuh.
Classes | |
struct | TempStorage |
The operations exposed by WarpReduce require a temporary memory allocation of this nested type for thread communication. This opaque storage can be allocated directly using the __shared__ keyword. Alternatively, it can be aliased to externally allocated memory (shared or global) or union 'd with other storage allocation types to facilitate memory reuse. More... | |
Public Methods | |
Collective constructors | |
__device__ __forceinline__ | WarpReduce (TempStorage &temp_storage) |
Collective constructor using the specified memory allocation as temporary storage. Logical warp and lane identifiers are constructed from threadIdx.x . More... | |
Summation reductions | |
__device__ __forceinline__ T | Sum (T input) |
Computes a warp-wide sum in the calling warp. The output is valid in warp lane0. More... | |
__device__ __forceinline__ T | Sum (T input, int valid_items) |
Computes a partially-full warp-wide sum in the calling warp. The output is valid in warp lane0. More... | |
template<typename FlagT > | |
__device__ __forceinline__ T | HeadSegmentedSum (T input, FlagT head_flag) |
Computes a segmented sum in the calling warp where segments are defined by head-flags. The sum of each segment is returned to the first lane in that segment (which always includes lane0). More... | |
template<typename FlagT > | |
__device__ __forceinline__ T | TailSegmentedSum (T input, FlagT tail_flag) |
Computes a segmented sum in the calling warp where segments are defined by tail-flags. The sum of each segment is returned to the first lane in that segment (which always includes lane0). More... | |
Generic reductions | |
template<typename ReductionOp > | |
__device__ __forceinline__ T | Reduce (T input, ReductionOp reduction_op) |
Computes a warp-wide reduction in the calling warp using the specified binary reduction functor. The output is valid in warp lane0. More... | |
template<typename ReductionOp > | |
__device__ __forceinline__ T | Reduce (T input, ReductionOp reduction_op, int valid_items) |
Computes a partially-full warp-wide reduction in the calling warp using the specified binary reduction functor. The output is valid in warp lane0. More... | |
template<typename ReductionOp , typename FlagT > | |
__device__ __forceinline__ T | HeadSegmentedReduce (T input, FlagT head_flag, ReductionOp reduction_op) |
Computes a segmented reduction in the calling warp where segments are defined by head-flags. The reduction of each segment is returned to the first lane in that segment (which always includes lane0). More... | |
template<typename ReductionOp , typename FlagT > | |
__device__ __forceinline__ T | TailSegmentedReduce (T input, FlagT tail_flag, ReductionOp reduction_op) |
Computes a segmented reduction in the calling warp where segments are defined by tail-flags. The reduction of each segment is returned to the first lane in that segment (which always includes lane0). More... | |
|
inline |
Collective constructor using the specified memory allocation as temporary storage. Logical warp and lane identifiers are constructed from threadIdx.x
.
[in] | temp_storage | Reference to memory allocation having layout type TempStorage |
Definition at line 203 of file warp_reduce.cuh.
|
inline |
Computes a warp-wide sum in the calling warp. The output is valid in warp lane0.
A subsequent __syncthreads()
threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., temp_storage
) is to be reused or repurposed.
thread_data
across the block of threads is {0, 1, 2, 3, ..., 127}
. The corresponding output aggregate
in threads 0, 32, 64, and 96 will 496
, 1520
, 2544
, and 3568
, respectively (and is undefined in other threads). [in] | input | Calling thread's input |
Definition at line 251 of file warp_reduce.cuh.
|
inline |
Computes a partially-full warp-wide sum in the calling warp. The output is valid in warp lane0.
All threads across the calling warp must agree on the same value for valid_items
. Otherwise the result is undefined.
A subsequent __syncthreads()
threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., temp_storage
) is to be reused or repurposed.
d_data
is {0, 1, 2, 3, 4, ...
and valid_items
is 4
. The corresponding output aggregate
in thread0 is 6
(and is undefined in other threads). [in] | input | Calling thread's input |
[in] | valid_items | Total number of valid items in the calling thread's logical warp (may be less than LOGICAL_WARP_THREADS ) |
Definition at line 295 of file warp_reduce.cuh.
|
inline |
Computes a segmented sum in the calling warp where segments are defined by head-flags. The sum of each segment is returned to the first lane in that segment (which always includes lane0).
A subsequent __syncthreads()
threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., temp_storage
) is to be reused or repurposed.
thread_data
and head_flag
across the block of threads is {0, 1, 2, 3, ..., 31
and is {1, 0, 0, 0, 1, 0, 0, 0, ..., 1, 0, 0, 0
, respectively. The corresponding output aggregate
in threads 0, 4, 8, etc. will be 6
, 22
, 38
, etc. (and is undefined in other threads).ReductionOp | [inferred] Binary reduction operator type having member T operator()(const T &a, const T &b) |
[in] | input | Calling thread's input |
[in] | head_flag | Head flag denoting whether or not input is the start of a new segment |
Definition at line 344 of file warp_reduce.cuh.
|
inline |
Computes a segmented sum in the calling warp where segments are defined by tail-flags. The sum of each segment is returned to the first lane in that segment (which always includes lane0).
A subsequent __syncthreads()
threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., temp_storage
) is to be reused or repurposed.
thread_data
and tail_flag
across the block of threads is {0, 1, 2, 3, ..., 31
and is {0, 0, 0, 1, 0, 0, 0, 1, ..., 0, 0, 0, 1
, respectively. The corresponding output aggregate
in threads 0, 4, 8, etc. will be 6
, 22
, 38
, etc. (and is undefined in other threads).ReductionOp | [inferred] Binary reduction operator type having member T operator()(const T &a, const T &b) |
[in] | input | Calling thread's input |
[in] | tail_flag | Head flag denoting whether or not input is the start of a new segment |
Definition at line 391 of file warp_reduce.cuh.
|
inline |
Computes a warp-wide reduction in the calling warp using the specified binary reduction functor. The output is valid in warp lane0.
Supports non-commutative reduction operators
A subsequent __syncthreads()
threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., temp_storage
) is to be reused or repurposed.
thread_data
across the block of threads is {0, 1, 2, 3, ..., 127}
. The corresponding output aggregate
in threads 0, 32, 64, and 96 will 31
, 63
, 95
, and 127
, respectively (and is undefined in other threads).ReductionOp | [inferred] Binary reduction operator type having member T operator()(const T &a, const T &b) |
[in] | input | Calling thread's input |
[in] | reduction_op | Binary reduction operator |
Definition at line 445 of file warp_reduce.cuh.
|
inline |
Computes a partially-full warp-wide reduction in the calling warp using the specified binary reduction functor. The output is valid in warp lane0.
All threads across the calling warp must agree on the same value for valid_items
. Otherwise the result is undefined.
Supports non-commutative reduction operators
A subsequent __syncthreads()
threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., temp_storage
) is to be reused or repurposed.
d_data
is {0, 1, 2, 3, 4, ...
and valid_items
is 4
. The corresponding output aggregate
in thread0 is 3
(and is undefined in other threads).ReductionOp | [inferred] Binary reduction operator type having member T operator()(const T &a, const T &b) |
[in] | input | Calling thread's input |
[in] | reduction_op | Binary reduction operator |
[in] | valid_items | Total number of valid items in the calling thread's logical warp (may be less than LOGICAL_WARP_THREADS ) |
Definition at line 494 of file warp_reduce.cuh.
|
inline |
Computes a segmented reduction in the calling warp where segments are defined by head-flags. The reduction of each segment is returned to the first lane in that segment (which always includes lane0).
Supports non-commutative reduction operators
A subsequent __syncthreads()
threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., temp_storage
) is to be reused or repurposed.
thread_data
and head_flag
across the block of threads is {0, 1, 2, 3, ..., 31
and is {1, 0, 0, 0, 1, 0, 0, 0, ..., 1, 0, 0, 0
, respectively. The corresponding output aggregate
in threads 0, 4, 8, etc. will be 3
, 7
, 11
, etc. (and is undefined in other threads).ReductionOp | [inferred] Binary reduction operator type having member T operator()(const T &a, const T &b) |
[in] | input | Calling thread's input |
[in] | head_flag | Head flag denoting whether or not input is the start of a new segment |
[in] | reduction_op | Reduction operator |
Definition at line 545 of file warp_reduce.cuh.
|
inline |
Computes a segmented reduction in the calling warp where segments are defined by tail-flags. The reduction of each segment is returned to the first lane in that segment (which always includes lane0).
Supports non-commutative reduction operators
A subsequent __syncthreads()
threadblock barrier should be invoked after calling this method if the collective's temporary storage (e.g., temp_storage
) is to be reused or repurposed.
thread_data
and tail_flag
across the block of threads is {0, 1, 2, 3, ..., 31
and is {0, 0, 0, 1, 0, 0, 0, 1, ..., 0, 0, 0, 1
, respectively. The corresponding output aggregate
in threads 0, 4, 8, etc. will be 3
, 7
, 11
, etc. (and is undefined in other threads).ReductionOp | [inferred] Binary reduction operator type having member T operator()(const T &a, const T &b) |
[in] | input | Calling thread's input |
[in] | tail_flag | Tail flag denoting whether or not input is the end of the current segment |
[in] | reduction_op | Reduction operator |
Definition at line 596 of file warp_reduce.cuh.