Detailed description

template< typename T, int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
class cub::BlockShuffle< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >

The BlockShuffle class provides collective methods for shuffling data partitioned across a CUDA thread block.

Template Parameters

T	The data type to be exchanged.
BLOCK_DIM_X	The thread block length in threads along the X dimension
BLOCK_DIM_Y	[optional] The thread block length in threads along the Y dimension (default: 1)
BLOCK_DIM_Z	[optional] The thread block length in threads along the Z dimension (default: 1)
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)

Overview: It is commonplace for blocks of threads to rearrange data items between threads. The BlockShuffle abstraction allows threads to efficiently shift items either (a) up to their successor or (b) down to their predecessor.

Classes
struct	TempStorage
	The operations exposed by BlockShuffle 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__	BlockShuffle ()
	Collective constructor using a private static allocation of shared memory as temporary storage. More...

__device__ __forceinline__	BlockShuffle (TempStorage &temp_storage)
	Collective constructor using the specified memory allocation as temporary storage. More...

Shuffle movement
__device__ __forceinline__ void	Offset (T input, T &output, int distance=1)
	Each thread_i obtains the `input` provided by thread_i+distance. The offset `distance` may be negative. More...

__device__ __forceinline__ void	Rotate (T input, T &output, unsigned int distance=1)
	Each thread_i obtains the `input` provided by thread_i+distance. More...

template<int ITEMS_PER_THREAD>
__device__ __forceinline__ void	Up (T(&input)[ITEMS_PER_THREAD], T(&prev)[ITEMS_PER_THREAD])
	The thread block rotates its blocked arrangement of `input` items, shifting it up by one item. More...

template<int ITEMS_PER_THREAD>
__device__ __forceinline__ void	Up (T(&input)[ITEMS_PER_THREAD], T(&prev)[ITEMS_PER_THREAD], T &block_suffix)
	The thread block rotates its blocked arrangement of `input` items, shifting it up by one item. All threads receive the `input` provided by thread_{BLOCK_THREADS-1}. More...

template<int ITEMS_PER_THREAD>
__device__ __forceinline__ void	Down (T(&input)[ITEMS_PER_THREAD], T(&prev)[ITEMS_PER_THREAD])
	The thread block rotates its blocked arrangement of `input` items, shifting it down by one item. More...

template<int ITEMS_PER_THREAD>
__device__ __forceinline__ void	Down (T(&input)[ITEMS_PER_THREAD], T(&prev)[ITEMS_PER_THREAD], T &block_prefix)
	The thread block rotates its blocked arrangement of input items, shifting it down by one item. All threads receive `input`[0] provided by thread₀. More...

Constructor & Destructor Documentation

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

__device__ __forceinline__ cub::BlockShuffle< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::BlockShuffle ( )

inline

Collective constructor using a private static allocation of shared memory as temporary storage.

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

__device__ __forceinline__ cub::BlockShuffle< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::BlockShuffle ( TempStorage & temp_storage)

inline

Collective constructor using the specified memory allocation as temporary storage.

Parameters

[in] temp_storage Reference to memory allocation having layout type TempStorage

Member Function Documentation

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

__device__ __forceinline__ void cub::BlockShuffle< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::Offset	(	T	input,
		T &	output,
		int	distance = `1`
	)

inline

Each thread_i obtains the input provided by thread_i+distance. The offset distance may be negative.

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.

Parameters

[in]	input	The input item from the calling thread (thread_i)
[out]	output	The `input` item from the successor (or predecessor) thread thread_i+distance (may be aliased to `input`). This value is only updated for for thread_i when 0 <= (i + `distance`) < `BLOCK_THREADS-1`
[in]	distance	Offset distance (may be negative)

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

__device__ __forceinline__ void cub::BlockShuffle< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::Rotate	(	T	input,
		T &	output,
		unsigned int	distance = `1`
	)

inline

Each thread_i obtains the input provided by thread_i+distance.

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.

Parameters

[in]	input	The calling thread's input item
[out]	output	The `input` item from thread thread_{(i+distance>)%BLOCK_THREADS} (may be aliased to `input`). This value is not updated for thread_{BLOCK_THREADS-1}
[in]	distance	Offset distance (0 < `distance` < `BLOCK_THREADS`)

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

template<int ITEMS_PER_THREAD>

__device__ __forceinline__ void cub::BlockShuffle< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::Up	(	T(&)	input[ITEMS_PER_THREAD],
		T(&)	prev[ITEMS_PER_THREAD]
	)

inline

The thread block rotates its blocked arrangement of input items, shifting it up by one item.

Assumes a blocked arrangement of (block-threads*items-per-thread) items across the thread block, where thread_i owns the i^th range of items-per-thread contiguous items. For multi-dimensional thread blocks, a row-major thread ordering is assumed.
Efficiency is increased with increased granularity ITEMS_PER_THREAD. Performance is also typically increased until the additional register pressure or shared memory allocation size causes SM occupancy to fall too low. Consider variants of cub::BlockLoad for efficiently gathering a blocked arrangement of elements across threads.
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.

Parameters

[in]	input	The calling thread's input items
[out]	prev	The corresponding predecessor items (may be aliased to `input`). The item `prev`[0] is not updated for thread₀.

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

template<int ITEMS_PER_THREAD>

__device__ __forceinline__ void cub::BlockShuffle< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::Up	(	T(&)	input[ITEMS_PER_THREAD],
		T(&)	prev[ITEMS_PER_THREAD],
		T &	block_suffix
	)

inline

The thread block rotates its blocked arrangement of input items, shifting it up by one item. All threads receive the input provided by thread_{BLOCK_THREADS-1}.

Assumes a blocked arrangement of (block-threads*items-per-thread) items across the thread block, where thread_i owns the i^th range of items-per-thread contiguous items. For multi-dimensional thread blocks, a row-major thread ordering is assumed.
Efficiency is increased with increased granularity ITEMS_PER_THREAD. Performance is also typically increased until the additional register pressure or shared memory allocation size causes SM occupancy to fall too low. Consider variants of cub::BlockLoad for efficiently gathering a blocked arrangement of elements across threads.
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.

Parameters

[in]	input	The calling thread's input items
[out]	prev	The corresponding predecessor items (may be aliased to `input`). The item `prev`[0] is not updated for thread₀.
[out]	block_suffix	The item `input`[ITEMS_PER_THREAD-1] from thread_{BLOCK_THREADS-1}, provided to all threads

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

template<int ITEMS_PER_THREAD>

__device__ __forceinline__ void cub::BlockShuffle< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::Down	(	T(&)	input[ITEMS_PER_THREAD],
		T(&)	prev[ITEMS_PER_THREAD]
	)

inline

The thread block rotates its blocked arrangement of input items, shifting it down by one item.

Assumes a blocked arrangement of (block-threads*items-per-thread) items across the thread block, where thread_i owns the i^th range of items-per-thread contiguous items. For multi-dimensional thread blocks, a row-major thread ordering is assumed.
Efficiency is increased with increased granularity ITEMS_PER_THREAD. Performance is also typically increased until the additional register pressure or shared memory allocation size causes SM occupancy to fall too low. Consider variants of cub::BlockLoad for efficiently gathering a blocked arrangement of elements across threads.
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.

Parameters

[in]	input	The calling thread's input items
[out]	prev	The corresponding predecessor items (may be aliased to `input`). The value `prev`[0] is not updated for thread_{BLOCK_THREADS-1}.

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

template<int ITEMS_PER_THREAD>

__device__ __forceinline__ void cub::BlockShuffle< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::Down	(	T(&)	input[ITEMS_PER_THREAD],
		T(&)	prev[ITEMS_PER_THREAD],
		T &	block_prefix
	)

inline

The thread block rotates its blocked arrangement of input items, shifting it down by one item. All threads receive input[0] provided by thread₀.

Assumes a blocked arrangement of (block-threads*items-per-thread) items across the thread block, where thread_i owns the i^th range of items-per-thread contiguous items. For multi-dimensional thread blocks, a row-major thread ordering is assumed.
Efficiency is increased with increased granularity ITEMS_PER_THREAD. Performance is also typically increased until the additional register pressure or shared memory allocation size causes SM occupancy to fall too low. Consider variants of cub::BlockLoad for efficiently gathering a blocked arrangement of elements across threads.
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.

Parameters

[in]	input	The calling thread's input items
[out]	prev	The corresponding predecessor items (may be aliased to `input`). The value `prev`[0] is not updated for thread_{BLOCK_THREADS-1}.
[out]	block_prefix	The item `input`[0] from thread₀, provided to all threads

The documentation for this class was generated from the following file:

block_shuffle.cuh

Detailed description

template< typename T, int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH> class cub::BlockShuffle< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >

Classes

Public Methods

Constructor & Destructor Documentation

Member Function Documentation

template< typename T, int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
class cub::BlockShuffle< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >