CUB  
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Macros Groups Pages
Classes | List of all members
cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH > Class Template Reference
Block-wide

Detailed description

template< typename T, int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
class cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >

The BlockScan class provides collective methods for computing a parallel prefix sum/scan of items partitioned across a CUDA thread block.

block_scan_logo.png
.
Template Parameters
TData type being scanned
BLOCK_DIM_XThe thread block length in threads along the X dimension
ALGORITHM[optional] cub::BlockScanAlgorithm enumerator specifying the underlying algorithm to use (default: cub::BLOCK_SCAN_RAKING)
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
  • Given a list of input elements and a binary reduction operator, a prefix scan produces an output list where each element is computed to be the reduction of the elements occurring earlier in the input list. Prefix sum connotes a prefix scan with the addition operator. The term inclusive indicates that the ith output reduction incorporates the ith input. The term exclusive indicates the ith input is not incorporated into the ith output reduction.
  • For multi-dimensional blocks, threads are linearly ranked in row-major order.
  • BlockScan can be optionally specialized by algorithm to accommodate different workload profiles:
    1. cub::BLOCK_SCAN_RAKING. An efficient (high throughput) "raking reduce-then-scan" prefix scan algorithm. More...
    2. cub::BLOCK_SCAN_RAKING_MEMOIZE. Similar to cub::BLOCK_SCAN_RAKING, but having higher throughput at the expense of additional register pressure for intermediate storage. More...
    3. cub::BLOCK_SCAN_WARP_SCANS. A quick (low latency) "tiled warpscans" prefix scan algorithm. More...
Performance Considerations
  • 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.
  • Uses special instructions when applicable (e.g., warp SHFL)
  • Uses synchronization-free communication between warp lanes when applicable
  • Invokes a minimal number of minimal block-wide synchronization barriers (only one or two depending on algorithm selection)
  • Incurs zero bank conflicts for most types
  • Computation is slightly more efficient (i.e., having lower instruction overhead) for:
    • Prefix sum variants (vs. generic scan)
    • The number of threads in the block is a multiple of the architecture's warp size
  • See cub::BlockScanAlgorithm for performance details regarding algorithmic alternatives
A Simple Example
Every thread in the block uses the BlockScan class by first specializing the BlockScan type, then instantiating an instance with parameters for communication, and finally invoking one or more collective member functions.
The code snippet below illustrates an exclusive prefix sum of 512 integer items that are partitioned in a blocked arrangement across 128 threads where each thread owns 4 consecutive items.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
__global__ void ExampleKernel(...)
{
// Specialize BlockScan for a 1D block of 128 threads of type int
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Obtain a segment of consecutive items that are blocked across threads
int thread_data[4];
...
// Collectively compute the block-wide exclusive prefix sum
BlockScan(temp_storage).ExclusiveSum(thread_data, thread_data);
Suppose the set of input thread_data across the block of threads is {[1,1,1,1], [1,1,1,1], ..., [1,1,1,1]}. The corresponding output thread_data in those threads will be {[0,1,2,3], [4,5,6,7], ..., [508,509,510,511]}.
Re-using dynamically allocating shared memory
The following example under the examples/block folder illustrates usage of dynamically shared memory with BlockReduce and how to re-purpose the same memory region: example_block_reduce_dyn_smem.cu

This example can be easily adapted to the storage required by BlockScan.

Examples:
example_block_scan.cu.

Classes

struct  TempStorage
 The operations exposed by BlockScan 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__ BlockScan ()
 Collective constructor using a private static allocation of shared memory as temporary storage. More...
 
__device__ __forceinline__ BlockScan (TempStorage &temp_storage)
 Collective constructor using the specified memory allocation as temporary storage. More...
 
Exclusive prefix sum operations
__device__ __forceinline__ void ExclusiveSum (T input, T &output)
 Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. The value of 0 is applied as the initial value, and is assigned to output in thread0. More...
 
__device__ __forceinline__ void ExclusiveSum (T input, T &output, T &block_aggregate)
 Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. The value of 0 is applied as the initial value, and is assigned to output in thread0. Also provides every thread with the block-wide block_aggregate of all inputs. More...
 
template<typename BlockPrefixCallbackOp >
__device__ __forceinline__ void ExclusiveSum (T input, T &output, BlockPrefixCallbackOp &block_prefix_callback_op)
 Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Instead of using 0 as the block-wide prefix, the call-back functor block_prefix_callback_op is invoked by the first warp in the block, and the value returned by lane0 in that warp is used as the "seed" value that logically prefixes the thread block's scan inputs. Also provides every thread with the block-wide block_aggregate of all inputs. More...
 
Exclusive prefix sum operations (multiple data per thread)
template<int ITEMS_PER_THREAD>
__device__ __forceinline__ void ExclusiveSum (T(&input)[ITEMS_PER_THREAD], T(&output)[ITEMS_PER_THREAD])
 Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. The value of 0 is applied as the initial value, and is assigned to output[0] in thread0. More...
 
template<int ITEMS_PER_THREAD>
__device__ __forceinline__ void ExclusiveSum (T(&input)[ITEMS_PER_THREAD], T(&output)[ITEMS_PER_THREAD], T &block_aggregate)
 Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. The value of 0 is applied as the initial value, and is assigned to output[0] in thread0. Also provides every thread with the block-wide block_aggregate of all inputs. More...
 
template<int ITEMS_PER_THREAD, typename BlockPrefixCallbackOp >
__device__ __forceinline__ void ExclusiveSum (T(&input)[ITEMS_PER_THREAD], T(&output)[ITEMS_PER_THREAD], BlockPrefixCallbackOp &block_prefix_callback_op)
 Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. Instead of using 0 as the block-wide prefix, the call-back functor block_prefix_callback_op is invoked by the first warp in the block, and the value returned by lane0 in that warp is used as the "seed" value that logically prefixes the thread block's scan inputs. Also provides every thread with the block-wide block_aggregate of all inputs. More...
 
Exclusive prefix scan operations
template<typename ScanOp >
__device__ __forceinline__ void ExclusiveScan (T input, T &output, T initial_value, ScanOp scan_op)
 Computes an exclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes one input element. More...
 
template<typename ScanOp >
__device__ __forceinline__ void ExclusiveScan (T input, T &output, T initial_value, ScanOp scan_op, T &block_aggregate)
 Computes an exclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes one input element. Also provides every thread with the block-wide block_aggregate of all inputs. More...
 
template<typename ScanOp , typename BlockPrefixCallbackOp >
__device__ __forceinline__ void ExclusiveScan (T input, T &output, ScanOp scan_op, BlockPrefixCallbackOp &block_prefix_callback_op)
 Computes an exclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes one input element. the call-back functor block_prefix_callback_op is invoked by the first warp in the block, and the value returned by lane0 in that warp is used as the "seed" value that logically prefixes the thread block's scan inputs. Also provides every thread with the block-wide block_aggregate of all inputs. More...
 
Exclusive prefix scan operations (multiple data per thread)
template<int ITEMS_PER_THREAD, typename ScanOp >
__device__ __forceinline__ void ExclusiveScan (T(&input)[ITEMS_PER_THREAD], T(&output)[ITEMS_PER_THREAD], T initial_value, ScanOp scan_op)
 Computes an exclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes an array of consecutive input elements. More...
 
template<int ITEMS_PER_THREAD, typename ScanOp >
__device__ __forceinline__ void ExclusiveScan (T(&input)[ITEMS_PER_THREAD], T(&output)[ITEMS_PER_THREAD], T initial_value, ScanOp scan_op, T &block_aggregate)
 Computes an exclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes an array of consecutive input elements. Also provides every thread with the block-wide block_aggregate of all inputs. More...
 
template<int ITEMS_PER_THREAD, typename ScanOp , typename BlockPrefixCallbackOp >
__device__ __forceinline__ void ExclusiveScan (T(&input)[ITEMS_PER_THREAD], T(&output)[ITEMS_PER_THREAD], ScanOp scan_op, BlockPrefixCallbackOp &block_prefix_callback_op)
 Computes an exclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes an array of consecutive input elements. the call-back functor block_prefix_callback_op is invoked by the first warp in the block, and the value returned by lane0 in that warp is used as the "seed" value that logically prefixes the thread block's scan inputs. Also provides every thread with the block-wide block_aggregate of all inputs. More...
 
Inclusive prefix sum operations
__device__ __forceinline__ void InclusiveSum (T input, T &output)
 Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. More...
 
__device__ __forceinline__ void InclusiveSum (T input, T &output, T &block_aggregate)
 Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Also provides every thread with the block-wide block_aggregate of all inputs. More...
 
template<typename BlockPrefixCallbackOp >
__device__ __forceinline__ void InclusiveSum (T input, T &output, BlockPrefixCallbackOp &block_prefix_callback_op)
 Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Instead of using 0 as the block-wide prefix, the call-back functor block_prefix_callback_op is invoked by the first warp in the block, and the value returned by lane0 in that warp is used as the "seed" value that logically prefixes the thread block's scan inputs. Also provides every thread with the block-wide block_aggregate of all inputs. More...
 
Inclusive prefix sum operations (multiple data per thread)
template<int ITEMS_PER_THREAD>
__device__ __forceinline__ void InclusiveSum (T(&input)[ITEMS_PER_THREAD], T(&output)[ITEMS_PER_THREAD])
 Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. More...
 
template<int ITEMS_PER_THREAD>
__device__ __forceinline__ void InclusiveSum (T(&input)[ITEMS_PER_THREAD], T(&output)[ITEMS_PER_THREAD], T &block_aggregate)
 Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. Also provides every thread with the block-wide block_aggregate of all inputs. More...
 
template<int ITEMS_PER_THREAD, typename BlockPrefixCallbackOp >
__device__ __forceinline__ void InclusiveSum (T(&input)[ITEMS_PER_THREAD], T(&output)[ITEMS_PER_THREAD], BlockPrefixCallbackOp &block_prefix_callback_op)
 Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. Instead of using 0 as the block-wide prefix, the call-back functor block_prefix_callback_op is invoked by the first warp in the block, and the value returned by lane0 in that warp is used as the "seed" value that logically prefixes the thread block's scan inputs. Also provides every thread with the block-wide block_aggregate of all inputs. More...
 
Inclusive prefix scan operations
template<typename ScanOp >
__device__ __forceinline__ void InclusiveScan (T input, T &output, ScanOp scan_op)
 Computes an inclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes one input element. More...
 
template<typename ScanOp >
__device__ __forceinline__ void InclusiveScan (T input, T &output, ScanOp scan_op, T &block_aggregate)
 Computes an inclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes one input element. Also provides every thread with the block-wide block_aggregate of all inputs. More...
 
template<typename ScanOp , typename BlockPrefixCallbackOp >
__device__ __forceinline__ void InclusiveScan (T input, T &output, ScanOp scan_op, BlockPrefixCallbackOp &block_prefix_callback_op)
 Computes an inclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes one input element. the call-back functor block_prefix_callback_op is invoked by the first warp in the block, and the value returned by lane0 in that warp is used as the "seed" value that logically prefixes the thread block's scan inputs. Also provides every thread with the block-wide block_aggregate of all inputs. More...
 
Inclusive prefix scan operations (multiple data per thread)
template<int ITEMS_PER_THREAD, typename ScanOp >
__device__ __forceinline__ void InclusiveScan (T(&input)[ITEMS_PER_THREAD], T(&output)[ITEMS_PER_THREAD], ScanOp scan_op)
 Computes an inclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes an array of consecutive input elements. More...
 
template<int ITEMS_PER_THREAD, typename ScanOp >
__device__ __forceinline__ void InclusiveScan (T(&input)[ITEMS_PER_THREAD], T(&output)[ITEMS_PER_THREAD], ScanOp scan_op, T &block_aggregate)
 Computes an inclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes an array of consecutive input elements. Also provides every thread with the block-wide block_aggregate of all inputs. More...
 
template<int ITEMS_PER_THREAD, typename ScanOp , typename BlockPrefixCallbackOp >
__device__ __forceinline__ void InclusiveScan (T(&input)[ITEMS_PER_THREAD], T(&output)[ITEMS_PER_THREAD], ScanOp scan_op, BlockPrefixCallbackOp &block_prefix_callback_op)
 Computes an inclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes an array of consecutive input elements. the call-back functor block_prefix_callback_op is invoked by the first warp in the block, and the value returned by lane0 in that warp is used as the "seed" value that logically prefixes the thread block's scan inputs. Also provides every thread with the block-wide block_aggregate of all inputs. More...
 

Constructor & Destructor Documentation

template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
__device__ __forceinline__ cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::BlockScan ( )
inline

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

template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
__device__ __forceinline__ cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::BlockScan ( TempStorage &  temp_storage)
inline

Collective constructor using the specified memory allocation as temporary storage.

Parameters
[in]temp_storageReference to memory allocation having layout type TempStorage

Member Function Documentation

template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
__device__ __forceinline__ void cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::ExclusiveSum ( T  input,
T &  output 
)
inline

Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. The value of 0 is applied as the initial value, and is assigned to output in thread0.

  • This operation assumes the value of obtained by the T's default constructor (or by zero-initialization if no user-defined default constructor exists) is suitable as the identity value zero for addition.
  • For multi-dimensional blocks, threads are linearly ranked in row-major order.
  • 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.
Snippet
The code snippet below illustrates an exclusive prefix sum of 128 integer items that are partitioned across 128 threads.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
__global__ void ExampleKernel(...)
{
// Specialize BlockScan for a 1D block of 128 threads of type int
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Obtain input item for each thread
int thread_data;
...
// Collectively compute the block-wide exclusive prefix sum
BlockScan(temp_storage).ExclusiveSum(thread_data, thread_data);
Suppose the set of input thread_data across the block of threads is 1, 1, ..., 1. The corresponding output thread_data in those threads will be 0, 1, ..., 127.
Parameters
[in]inputCalling thread's input item
[out]outputCalling thread's output item (may be aliased to input)
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
__device__ __forceinline__ void cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::ExclusiveSum ( T  input,
T &  output,
T &  block_aggregate 
)
inline

Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. The value of 0 is applied as the initial value, and is assigned to output in thread0. Also provides every thread with the block-wide block_aggregate of all inputs.

  • This operation assumes the value of obtained by the T's default constructor (or by zero-initialization if no user-defined default constructor exists) is suitable as the identity value zero for addition.
  • For multi-dimensional blocks, threads are linearly ranked in row-major order.
  • 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.
Snippet
The code snippet below illustrates an exclusive prefix sum of 128 integer items that are partitioned across 128 threads.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
__global__ void ExampleKernel(...)
{
// Specialize BlockScan for a 1D block of 128 threads of type int
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Obtain input item for each thread
int thread_data;
...
// Collectively compute the block-wide exclusive prefix sum
int block_aggregate;
BlockScan(temp_storage).ExclusiveSum(thread_data, thread_data, block_aggregate);
Suppose the set of input thread_data across the block of threads is 1, 1, ..., 1. The corresponding output thread_data in those threads will be 0, 1, ..., 127. Furthermore the value 128 will be stored in block_aggregate for all threads.
Parameters
[in]inputCalling thread's input item
[out]outputCalling thread's output item (may be aliased to input)
[out]block_aggregateblock-wide aggregate reduction of input items
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
template<typename BlockPrefixCallbackOp >
__device__ __forceinline__ void cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::ExclusiveSum ( T  input,
T &  output,
BlockPrefixCallbackOp &  block_prefix_callback_op 
)
inline

Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Instead of using 0 as the block-wide prefix, the call-back functor block_prefix_callback_op is invoked by the first warp in the block, and the value returned by lane0 in that warp is used as the "seed" value that logically prefixes the thread block's scan inputs. Also provides every thread with the block-wide block_aggregate of all inputs.

  • This operation assumes the value of obtained by the T's default constructor (or by zero-initialization if no user-defined default constructor exists) is suitable as the identity value zero for addition.
  • The block_prefix_callback_op functor must implement a member function T operator()(T block_aggregate). The functor's input parameter block_aggregate is the same value also returned by the scan operation. The functor will be invoked by the first warp of threads in the block, however only the return value from lane0 is applied as the block-wide prefix. Can be stateful.
  • For multi-dimensional blocks, threads are linearly ranked in row-major order.
  • 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.
Snippet
The code snippet below illustrates a single thread block that progressively computes an exclusive prefix sum over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 128 integer items that are partitioned across 128 threads.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
// A stateful callback functor that maintains a running prefix to be applied
// during consecutive scan operations.
struct BlockPrefixCallbackOp
{
// Running prefix
int running_total;
// Constructor
__device__ BlockPrefixCallbackOp(int running_total) : running_total(running_total) {}
// Callback operator to be entered by the first warp of threads in the block.
// Thread-0 is responsible for returning a value for seeding the block-wide scan.
__device__ int operator()(int block_aggregate)
{
int old_prefix = running_total;
running_total += block_aggregate;
return old_prefix;
}
};
__global__ void ExampleKernel(int *d_data, int num_items, ...)
{
// Specialize BlockScan for a 1D block of 128 threads
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Initialize running total
BlockPrefixCallbackOp prefix_op(0);
// Have the block iterate over segments of items
for (int block_offset = 0; block_offset < num_items; block_offset += 128)
{
// Load a segment of consecutive items that are blocked across threads
int thread_data = d_data[block_offset];
// Collectively compute the block-wide exclusive prefix sum
BlockScan(temp_storage).ExclusiveSum(
thread_data, thread_data, prefix_op);
CTA_SYNC();
// Store scanned items to output segment
d_data[block_offset] = thread_data;
}
Suppose the input d_data is 1, 1, 1, 1, 1, 1, 1, 1, .... The corresponding output for the first segment will be 0, 1, ..., 127. The output for the second segment will be 128, 129, ..., 255.
Template Parameters
BlockPrefixCallbackOp[inferred] Call-back functor type having member T operator()(T block_aggregate)
Parameters
[in]inputCalling thread's input item
[out]outputCalling thread's output item (may be aliased to input)
[in,out]block_prefix_callback_op[warp0 only] Call-back functor for specifying a block-wide prefix to be applied to the logical input sequence.
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, 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::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::ExclusiveSum ( T(&)  input[ITEMS_PER_THREAD],
T(&)  output[ITEMS_PER_THREAD] 
)
inline

Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. The value of 0 is applied as the initial value, and is assigned to output[0] in thread0.

  • This operation assumes the value of obtained by the T's default constructor (or by zero-initialization if no user-defined default constructor exists) is suitable as the identity value zero for addition.
  • Assumes a blocked arrangement of (block-threads*items-per-thread) items across the thread block, where threadi owns the ith 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.
Snippet
The code snippet below illustrates an exclusive prefix sum of 512 integer items that are partitioned in a blocked arrangement across 128 threads where each thread owns 4 consecutive items.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
__global__ void ExampleKernel(...)
{
// Specialize BlockScan for a 1D block of 128 threads of type int
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Obtain a segment of consecutive items that are blocked across threads
int thread_data[4];
...
// Collectively compute the block-wide exclusive prefix sum
BlockScan(temp_storage).ExclusiveSum(thread_data, thread_data);
Suppose the set of input thread_data across the block of threads is { [1,1,1,1], [1,1,1,1], ..., [1,1,1,1] }. The corresponding output thread_data in those threads will be { [0,1,2,3], [4,5,6,7], ..., [508,509,510,511] }.
Template Parameters
ITEMS_PER_THREAD[inferred] The number of consecutive items partitioned onto each thread.
Parameters
[in]inputCalling thread's input items
[out]outputCalling thread's output items (may be aliased to input)
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, 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::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::ExclusiveSum ( T(&)  input[ITEMS_PER_THREAD],
T(&)  output[ITEMS_PER_THREAD],
T &  block_aggregate 
)
inline

Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. The value of 0 is applied as the initial value, and is assigned to output[0] in thread0. Also provides every thread with the block-wide block_aggregate of all inputs.

  • This operation assumes the value of obtained by the T's default constructor (or by zero-initialization if no user-defined default constructor exists) is suitable as the identity value zero for addition.
  • Assumes a blocked arrangement of (block-threads*items-per-thread) items across the thread block, where threadi owns the ith 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.
Snippet
The code snippet below illustrates an exclusive prefix sum of 512 integer items that are partitioned in a blocked arrangement across 128 threads where each thread owns 4 consecutive items.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
__global__ void ExampleKernel(...)
{
// Specialize BlockScan for a 1D block of 128 threads of type int
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Obtain a segment of consecutive items that are blocked across threads
int thread_data[4];
...
// Collectively compute the block-wide exclusive prefix sum
int block_aggregate;
BlockScan(temp_storage).ExclusiveSum(thread_data, thread_data, block_aggregate);
Suppose the set of input thread_data across the block of threads is { [1,1,1,1], [1,1,1,1], ..., [1,1,1,1] }. The corresponding output thread_data in those threads will be { [0,1,2,3], [4,5,6,7], ..., [508,509,510,511] }. Furthermore the value 512 will be stored in block_aggregate for all threads.
Template Parameters
ITEMS_PER_THREAD[inferred] The number of consecutive items partitioned onto each thread.
Parameters
[in]inputCalling thread's input items
[out]outputCalling thread's output items (may be aliased to input)
[out]block_aggregateblock-wide aggregate reduction of input items
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
template<int ITEMS_PER_THREAD, typename BlockPrefixCallbackOp >
__device__ __forceinline__ void cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::ExclusiveSum ( T(&)  input[ITEMS_PER_THREAD],
T(&)  output[ITEMS_PER_THREAD],
BlockPrefixCallbackOp &  block_prefix_callback_op 
)
inline

Computes an exclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. Instead of using 0 as the block-wide prefix, the call-back functor block_prefix_callback_op is invoked by the first warp in the block, and the value returned by lane0 in that warp is used as the "seed" value that logically prefixes the thread block's scan inputs. Also provides every thread with the block-wide block_aggregate of all inputs.

  • This operation assumes the value of obtained by the T's default constructor (or by zero-initialization if no user-defined default constructor exists) is suitable as the identity value zero for addition.
  • The block_prefix_callback_op functor must implement a member function T operator()(T block_aggregate). The functor's input parameter block_aggregate is the same value also returned by the scan operation. The functor will be invoked by the first warp of threads in the block, however only the return value from lane0 is applied as the block-wide prefix. Can be stateful.
  • Assumes a blocked arrangement of (block-threads*items-per-thread) items across the thread block, where threadi owns the ith 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.
Snippet
The code snippet below illustrates a single thread block that progressively computes an exclusive prefix sum over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 512 integer items that are partitioned in a blocked arrangement across 128 threads where each thread owns 4 consecutive items.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
// A stateful callback functor that maintains a running prefix to be applied
// during consecutive scan operations.
struct BlockPrefixCallbackOp
{
// Running prefix
int running_total;
// Constructor
__device__ BlockPrefixCallbackOp(int running_total) : running_total(running_total) {}
// Callback operator to be entered by the first warp of threads in the block.
// Thread-0 is responsible for returning a value for seeding the block-wide scan.
__device__ int operator()(int block_aggregate)
{
int old_prefix = running_total;
running_total += block_aggregate;
return old_prefix;
}
};
__global__ void ExampleKernel(int *d_data, int num_items, ...)
{
// Specialize BlockLoad, BlockStore, and BlockScan for a 1D block of 128 threads, 4 ints per thread
typedef cub::BlockLoad<int*, 128, 4, BLOCK_LOAD_TRANSPOSE> BlockLoad;
typedef cub::BlockStore<int, 128, 4, BLOCK_STORE_TRANSPOSE> BlockStore;
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate aliased shared memory for BlockLoad, BlockStore, and BlockScan
__shared__ union {
typename BlockLoad::TempStorage load;
typename BlockScan::TempStorage scan;
typename BlockStore::TempStorage store;
} temp_storage;
// Initialize running total
BlockPrefixCallbackOp prefix_op(0);
// Have the block iterate over segments of items
for (int block_offset = 0; block_offset < num_items; block_offset += 128 * 4)
{
// Load a segment of consecutive items that are blocked across threads
int thread_data[4];
BlockLoad(temp_storage.load).Load(d_data + block_offset, thread_data);
CTA_SYNC();
// Collectively compute the block-wide exclusive prefix sum
int block_aggregate;
BlockScan(temp_storage.scan).ExclusiveSum(
thread_data, thread_data, prefix_op);
CTA_SYNC();
// Store scanned items to output segment
BlockStore(temp_storage.store).Store(d_data + block_offset, thread_data);
CTA_SYNC();
}
Suppose the input d_data is 1, 1, 1, 1, 1, 1, 1, 1, .... The corresponding output for the first segment will be 0, 1, 2, 3, ..., 510, 511. The output for the second segment will be 512, 513, 514, 515, ..., 1022, 1023.
Template Parameters
ITEMS_PER_THREAD[inferred] The number of consecutive items partitioned onto each thread.
BlockPrefixCallbackOp[inferred] Call-back functor type having member T operator()(T block_aggregate)
Parameters
[in]inputCalling thread's input items
[out]outputCalling thread's output items (may be aliased to input)
[in,out]block_prefix_callback_op[warp0 only] Call-back functor for specifying a block-wide prefix to be applied to the logical input sequence.
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
template<typename ScanOp >
__device__ __forceinline__ void cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::ExclusiveScan ( T  input,
T &  output,
T  initial_value,
ScanOp  scan_op 
)
inline

Computes an exclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes one input element.

  • Supports non-commutative scan operators.
  • For multi-dimensional blocks, threads are linearly ranked in row-major order.
  • 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.
Snippet
The code snippet below illustrates an exclusive prefix max scan of 128 integer items that are partitioned across 128 threads.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
__global__ void ExampleKernel(...)
{
// Specialize BlockScan for a 1D block of 128 threads of type int
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Obtain input item for each thread
int thread_data;
...
// Collectively compute the block-wide exclusive prefix max scan
BlockScan(temp_storage).ExclusiveScan(thread_data, thread_data, INT_MIN, cub::Max());
Suppose the set of input thread_data across the block of threads is 0, -1, 2, -3, ..., 126, -127. The corresponding output thread_data in those threads will be INT_MIN, 0, 0, 2, ..., 124, 126.
Template Parameters
ScanOp[inferred] Binary scan functor type having member T operator()(const T &a, const T &b)
Parameters
[in]inputCalling thread's input item
[out]outputCalling thread's output item (may be aliased to input)
[in]initial_valueInitial value to seed the exclusive scan (and is assigned to output[0] in thread0)
[in]scan_opBinary scan functor
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
template<typename ScanOp >
__device__ __forceinline__ void cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::ExclusiveScan ( T  input,
T &  output,
T  initial_value,
ScanOp  scan_op,
T &  block_aggregate 
)
inline

Computes an exclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes one input element. Also provides every thread with the block-wide block_aggregate of all inputs.

  • Supports non-commutative scan operators.
  • For multi-dimensional blocks, threads are linearly ranked in row-major order.
  • 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.
Snippet
The code snippet below illustrates an exclusive prefix max scan of 128 integer items that are partitioned across 128 threads.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
__global__ void ExampleKernel(...)
{
// Specialize BlockScan for a 1D block of 128 threads of type int
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Obtain input item for each thread
int thread_data;
...
// Collectively compute the block-wide exclusive prefix max scan
int block_aggregate;
BlockScan(temp_storage).ExclusiveScan(thread_data, thread_data, INT_MIN, cub::Max(), block_aggregate);
Suppose the set of input thread_data across the block of threads is 0, -1, 2, -3, ..., 126, -127. The corresponding output thread_data in those threads will be INT_MIN, 0, 0, 2, ..., 124, 126. Furthermore the value 126 will be stored in block_aggregate for all threads.
Template Parameters
ScanOp[inferred] Binary scan functor type having member T operator()(const T &a, const T &b)
Parameters
[in]inputCalling thread's input items
[out]outputCalling thread's output items (may be aliased to input)
[in]initial_valueInitial value to seed the exclusive scan (and is assigned to output[0] in thread0)
[in]scan_opBinary scan functor
[out]block_aggregateblock-wide aggregate reduction of input items
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
template<typename ScanOp , typename BlockPrefixCallbackOp >
__device__ __forceinline__ void cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::ExclusiveScan ( T  input,
T &  output,
ScanOp  scan_op,
BlockPrefixCallbackOp &  block_prefix_callback_op 
)
inline

Computes an exclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes one input element. the call-back functor block_prefix_callback_op is invoked by the first warp in the block, and the value returned by lane0 in that warp is used as the "seed" value that logically prefixes the thread block's scan inputs. Also provides every thread with the block-wide block_aggregate of all inputs.

  • The block_prefix_callback_op functor must implement a member function T operator()(T block_aggregate). The functor's input parameter block_aggregate is the same value also returned by the scan operation. The functor will be invoked by the first warp of threads in the block, however only the return value from lane0 is applied as the block-wide prefix. Can be stateful.
  • Supports non-commutative scan operators.
  • For multi-dimensional blocks, threads are linearly ranked in row-major order.
  • 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.
Snippet
The code snippet below illustrates a single thread block that progressively computes an exclusive prefix max scan over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 128 integer items that are partitioned across 128 threads.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
// A stateful callback functor that maintains a running prefix to be applied
// during consecutive scan operations.
struct BlockPrefixCallbackOp
{
// Running prefix
int running_total;
// Constructor
__device__ BlockPrefixCallbackOp(int running_total) : running_total(running_total) {}
// Callback operator to be entered by the first warp of threads in the block.
// Thread-0 is responsible for returning a value for seeding the block-wide scan.
__device__ int operator()(int block_aggregate)
{
int old_prefix = running_total;
running_total = (block_aggregate > old_prefix) ? block_aggregate : old_prefix;
return old_prefix;
}
};
__global__ void ExampleKernel(int *d_data, int num_items, ...)
{
// Specialize BlockScan for a 1D block of 128 threads
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Initialize running total
BlockPrefixCallbackOp prefix_op(INT_MIN);
// Have the block iterate over segments of items
for (int block_offset = 0; block_offset < num_items; block_offset += 128)
{
// Load a segment of consecutive items that are blocked across threads
int thread_data = d_data[block_offset];
// Collectively compute the block-wide exclusive prefix max scan
BlockScan(temp_storage).ExclusiveScan(
thread_data, thread_data, INT_MIN, cub::Max(), prefix_op);
CTA_SYNC();
// Store scanned items to output segment
d_data[block_offset] = thread_data;
}
Suppose the input d_data is 0, -1, 2, -3, 4, -5, .... The corresponding output for the first segment will be INT_MIN, 0, 0, 2, ..., 124, 126. The output for the second segment will be 126, 128, 128, 130, ..., 252, 254.
Template Parameters
ScanOp[inferred] Binary scan functor type having member T operator()(const T &a, const T &b)
BlockPrefixCallbackOp[inferred] Call-back functor type having member T operator()(T block_aggregate)
Parameters
[in]inputCalling thread's input item
[out]outputCalling thread's output item (may be aliased to input)
[in]scan_opBinary scan functor
[in,out]block_prefix_callback_op[warp0 only] Call-back functor for specifying a block-wide prefix to be applied to the logical input sequence.
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
template<int ITEMS_PER_THREAD, typename ScanOp >
__device__ __forceinline__ void cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::ExclusiveScan ( T(&)  input[ITEMS_PER_THREAD],
T(&)  output[ITEMS_PER_THREAD],
T  initial_value,
ScanOp  scan_op 
)
inline

Computes an exclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes an array of consecutive input elements.

  • Supports non-commutative scan operators.
  • Assumes a blocked arrangement of (block-threads*items-per-thread) items across the thread block, where threadi owns the ith 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.
Snippet
The code snippet below illustrates an exclusive prefix max scan of 512 integer items that are partitioned in a blocked arrangement across 128 threads where each thread owns 4 consecutive items.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
__global__ void ExampleKernel(...)
{
// Specialize BlockScan for a 1D block of 128 threads of type int
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Obtain a segment of consecutive items that are blocked across threads
int thread_data[4];
...
// Collectively compute the block-wide exclusive prefix max scan
BlockScan(temp_storage).ExclusiveScan(thread_data, thread_data, INT_MIN, cub::Max());
Suppose the set of input thread_data across the block of threads is { [0,-1,2,-3], [4,-5,6,-7], ..., [508,-509,510,-511] }. The corresponding output thread_data in those threads will be { [INT_MIN,0,0,2], [2,4,4,6], ..., [506,508,508,510] }.
Template Parameters
ITEMS_PER_THREAD[inferred] The number of consecutive items partitioned onto each thread.
ScanOp[inferred] Binary scan functor type having member T operator()(const T &a, const T &b)
Parameters
[in]inputCalling thread's input items
[out]outputCalling thread's output items (may be aliased to input)
[in]initial_valueInitial value to seed the exclusive scan (and is assigned to output[0] in thread0)
[in]scan_opBinary scan functor
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
template<int ITEMS_PER_THREAD, typename ScanOp >
__device__ __forceinline__ void cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::ExclusiveScan ( T(&)  input[ITEMS_PER_THREAD],
T(&)  output[ITEMS_PER_THREAD],
T  initial_value,
ScanOp  scan_op,
T &  block_aggregate 
)
inline

Computes an exclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes an array of consecutive input elements. Also provides every thread with the block-wide block_aggregate of all inputs.

  • Supports non-commutative scan operators.
  • Assumes a blocked arrangement of (block-threads*items-per-thread) items across the thread block, where threadi owns the ith 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.
Snippet
The code snippet below illustrates an exclusive prefix max scan of 512 integer items that are partitioned in a blocked arrangement across 128 threads where each thread owns 4 consecutive items.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
__global__ void ExampleKernel(...)
{
// Specialize BlockScan for a 1D block of 128 threads of type int
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Obtain a segment of consecutive items that are blocked across threads
int thread_data[4];
...
// Collectively compute the block-wide exclusive prefix max scan
int block_aggregate;
BlockScan(temp_storage).ExclusiveScan(thread_data, thread_data, INT_MIN, cub::Max(), block_aggregate);
Suppose the set of input thread_data across the block of threads is { [0,-1,2,-3], [4,-5,6,-7], ..., [508,-509,510,-511] }. The corresponding output thread_data in those threads will be { [INT_MIN,0,0,2], [2,4,4,6], ..., [506,508,508,510] }. Furthermore the value 510 will be stored in block_aggregate for all threads.
Template Parameters
ITEMS_PER_THREAD[inferred] The number of consecutive items partitioned onto each thread.
ScanOp[inferred] Binary scan functor type having member T operator()(const T &a, const T &b)
Parameters
[in]inputCalling thread's input items
[out]outputCalling thread's output items (may be aliased to input)
[in]initial_valueInitial value to seed the exclusive scan (and is assigned to output[0] in thread0)
[in]scan_opBinary scan functor
[out]block_aggregateblock-wide aggregate reduction of input items
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
template<int ITEMS_PER_THREAD, typename ScanOp , typename BlockPrefixCallbackOp >
__device__ __forceinline__ void cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::ExclusiveScan ( T(&)  input[ITEMS_PER_THREAD],
T(&)  output[ITEMS_PER_THREAD],
ScanOp  scan_op,
BlockPrefixCallbackOp &  block_prefix_callback_op 
)
inline

Computes an exclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes an array of consecutive input elements. the call-back functor block_prefix_callback_op is invoked by the first warp in the block, and the value returned by lane0 in that warp is used as the "seed" value that logically prefixes the thread block's scan inputs. Also provides every thread with the block-wide block_aggregate of all inputs.

  • The block_prefix_callback_op functor must implement a member function T operator()(T block_aggregate). The functor's input parameter block_aggregate is the same value also returned by the scan operation. The functor will be invoked by the first warp of threads in the block, however only the return value from lane0 is applied as the block-wide prefix. Can be stateful.
  • Supports non-commutative scan operators.
  • Assumes a blocked arrangement of (block-threads*items-per-thread) items across the thread block, where threadi owns the ith 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.
Snippet
The code snippet below illustrates a single thread block that progressively computes an exclusive prefix max scan over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 128 integer items that are partitioned across 128 threads.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
// A stateful callback functor that maintains a running prefix to be applied
// during consecutive scan operations.
struct BlockPrefixCallbackOp
{
// Running prefix
int running_total;
// Constructor
__device__ BlockPrefixCallbackOp(int running_total) : running_total(running_total) {}
// Callback operator to be entered by the first warp of threads in the block.
// Thread-0 is responsible for returning a value for seeding the block-wide scan.
__device__ int operator()(int block_aggregate)
{
int old_prefix = running_total;
running_total = (block_aggregate > old_prefix) ? block_aggregate : old_prefix;
return old_prefix;
}
};
__global__ void ExampleKernel(int *d_data, int num_items, ...)
{
// Specialize BlockLoad, BlockStore, and BlockScan for a 1D block of 128 threads, 4 ints per thread
typedef cub::BlockLoad<int*, 128, 4, BLOCK_LOAD_TRANSPOSE> BlockLoad;
typedef cub::BlockStore<int, 128, 4, BLOCK_STORE_TRANSPOSE> BlockStore;
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate aliased shared memory for BlockLoad, BlockStore, and BlockScan
__shared__ union {
typename BlockLoad::TempStorage load;
typename BlockScan::TempStorage scan;
typename BlockStore::TempStorage store;
} temp_storage;
// Initialize running total
BlockPrefixCallbackOp prefix_op(0);
// Have the block iterate over segments of items
for (int block_offset = 0; block_offset < num_items; block_offset += 128 * 4)
{
// Load a segment of consecutive items that are blocked across threads
int thread_data[4];
BlockLoad(temp_storage.load).Load(d_data + block_offset, thread_data);
CTA_SYNC();
// Collectively compute the block-wide exclusive prefix max scan
BlockScan(temp_storage.scan).ExclusiveScan(
thread_data, thread_data, INT_MIN, cub::Max(), prefix_op);
CTA_SYNC();
// Store scanned items to output segment
BlockStore(temp_storage.store).Store(d_data + block_offset, thread_data);
CTA_SYNC();
}
Suppose the input d_data is 0, -1, 2, -3, 4, -5, .... The corresponding output for the first segment will be INT_MIN, 0, 0, 2, 2, 4, ..., 508, 510. The output for the second segment will be 510, 512, 512, 514, 514, 516, ..., 1020, 1022.
Template Parameters
ITEMS_PER_THREAD[inferred] The number of consecutive items partitioned onto each thread.
ScanOp[inferred] Binary scan functor type having member T operator()(const T &a, const T &b)
BlockPrefixCallbackOp[inferred] Call-back functor type having member T operator()(T block_aggregate)
Parameters
[in]inputCalling thread's input items
[out]outputCalling thread's output items (may be aliased to input)
[in]scan_opBinary scan functor
[in,out]block_prefix_callback_op[warp0 only] Call-back functor for specifying a block-wide prefix to be applied to the logical input sequence.
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
__device__ __forceinline__ void cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::InclusiveSum ( T  input,
T &  output 
)
inline

Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element.

  • For multi-dimensional blocks, threads are linearly ranked in row-major order.
  • 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.
Snippet
The code snippet below illustrates an inclusive prefix sum of 128 integer items that are partitioned across 128 threads.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
__global__ void ExampleKernel(...)
{
// Specialize BlockScan for a 1D block of 128 threads of type int
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Obtain input item for each thread
int thread_data;
...
// Collectively compute the block-wide inclusive prefix sum
BlockScan(temp_storage).InclusiveSum(thread_data, thread_data);
Suppose the set of input thread_data across the block of threads is 1, 1, ..., 1. The corresponding output thread_data in those threads will be 1, 2, ..., 128.
Parameters
[in]inputCalling thread's input item
[out]outputCalling thread's output item (may be aliased to input)
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
__device__ __forceinline__ void cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::InclusiveSum ( T  input,
T &  output,
T &  block_aggregate 
)
inline

Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Also provides every thread with the block-wide block_aggregate of all inputs.

  • For multi-dimensional blocks, threads are linearly ranked in row-major order.
  • 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.
Snippet
The code snippet below illustrates an inclusive prefix sum of 128 integer items that are partitioned across 128 threads.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
__global__ void ExampleKernel(...)
{
// Specialize BlockScan for a 1D block of 128 threads of type int
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Obtain input item for each thread
int thread_data;
...
// Collectively compute the block-wide inclusive prefix sum
int block_aggregate;
BlockScan(temp_storage).InclusiveSum(thread_data, thread_data, block_aggregate);
Suppose the set of input thread_data across the block of threads is 1, 1, ..., 1. The corresponding output thread_data in those threads will be 1, 2, ..., 128. Furthermore the value 128 will be stored in block_aggregate for all threads.
Parameters
[in]inputCalling thread's input item
[out]outputCalling thread's output item (may be aliased to input)
[out]block_aggregateblock-wide aggregate reduction of input items
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
template<typename BlockPrefixCallbackOp >
__device__ __forceinline__ void cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::InclusiveSum ( T  input,
T &  output,
BlockPrefixCallbackOp &  block_prefix_callback_op 
)
inline

Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes one input element. Instead of using 0 as the block-wide prefix, the call-back functor block_prefix_callback_op is invoked by the first warp in the block, and the value returned by lane0 in that warp is used as the "seed" value that logically prefixes the thread block's scan inputs. Also provides every thread with the block-wide block_aggregate of all inputs.

  • The block_prefix_callback_op functor must implement a member function T operator()(T block_aggregate). The functor's input parameter block_aggregate is the same value also returned by the scan operation. The functor will be invoked by the first warp of threads in the block, however only the return value from lane0 is applied as the block-wide prefix. Can be stateful.
  • For multi-dimensional blocks, threads are linearly ranked in row-major order.
  • 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.
Snippet
The code snippet below illustrates a single thread block that progressively computes an inclusive prefix sum over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 128 integer items that are partitioned across 128 threads.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
// A stateful callback functor that maintains a running prefix to be applied
// during consecutive scan operations.
struct BlockPrefixCallbackOp
{
// Running prefix
int running_total;
// Constructor
__device__ BlockPrefixCallbackOp(int running_total) : running_total(running_total) {}
// Callback operator to be entered by the first warp of threads in the block.
// Thread-0 is responsible for returning a value for seeding the block-wide scan.
__device__ int operator()(int block_aggregate)
{
int old_prefix = running_total;
running_total += block_aggregate;
return old_prefix;
}
};
__global__ void ExampleKernel(int *d_data, int num_items, ...)
{
// Specialize BlockScan for a 1D block of 128 threads
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Initialize running total
BlockPrefixCallbackOp prefix_op(0);
// Have the block iterate over segments of items
for (int block_offset = 0; block_offset < num_items; block_offset += 128)
{
// Load a segment of consecutive items that are blocked across threads
int thread_data = d_data[block_offset];
// Collectively compute the block-wide inclusive prefix sum
BlockScan(temp_storage).InclusiveSum(
thread_data, thread_data, prefix_op);
CTA_SYNC();
// Store scanned items to output segment
d_data[block_offset] = thread_data;
}
Suppose the input d_data is 1, 1, 1, 1, 1, 1, 1, 1, .... The corresponding output for the first segment will be 1, 2, ..., 128. The output for the second segment will be 129, 130, ..., 256.
Template Parameters
BlockPrefixCallbackOp[inferred] Call-back functor type having member T operator()(T block_aggregate)
Parameters
[in]inputCalling thread's input item
[out]outputCalling thread's output item (may be aliased to input)
[in,out]block_prefix_callback_op[warp0 only] Call-back functor for specifying a block-wide prefix to be applied to the logical input sequence.
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, 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::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::InclusiveSum ( T(&)  input[ITEMS_PER_THREAD],
T(&)  output[ITEMS_PER_THREAD] 
)
inline

Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements.

  • Assumes a blocked arrangement of (block-threads*items-per-thread) items across the thread block, where threadi owns the ith 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.
Snippet
The code snippet below illustrates an inclusive prefix sum of 512 integer items that are partitioned in a blocked arrangement across 128 threads where each thread owns 4 consecutive items.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
__global__ void ExampleKernel(...)
{
// Specialize BlockScan for a 1D block of 128 threads of type int
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Obtain a segment of consecutive items that are blocked across threads
int thread_data[4];
...
// Collectively compute the block-wide inclusive prefix sum
BlockScan(temp_storage).InclusiveSum(thread_data, thread_data);
Suppose the set of input thread_data across the block of threads is { [1,1,1,1], [1,1,1,1], ..., [1,1,1,1] }. The corresponding output thread_data in those threads will be { [1,2,3,4], [5,6,7,8], ..., [509,510,511,512] }.
Template Parameters
ITEMS_PER_THREAD[inferred] The number of consecutive items partitioned onto each thread.
Parameters
[in]inputCalling thread's input items
[out]outputCalling thread's output items (may be aliased to input)
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, 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::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::InclusiveSum ( T(&)  input[ITEMS_PER_THREAD],
T(&)  output[ITEMS_PER_THREAD],
T &  block_aggregate 
)
inline

Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. Also provides every thread with the block-wide block_aggregate of all inputs.

  • Assumes a blocked arrangement of (block-threads*items-per-thread) items across the thread block, where threadi owns the ith 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.
Snippet
The code snippet below illustrates an inclusive prefix sum of 512 integer items that are partitioned in a blocked arrangement across 128 threads where each thread owns 4 consecutive items.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
__global__ void ExampleKernel(...)
{
// Specialize BlockScan for a 1D block of 128 threads of type int
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Obtain a segment of consecutive items that are blocked across threads
int thread_data[4];
...
// Collectively compute the block-wide inclusive prefix sum
int block_aggregate;
BlockScan(temp_storage).InclusiveSum(thread_data, thread_data, block_aggregate);
Suppose the set of input thread_data across the block of threads is { [1,1,1,1], [1,1,1,1], ..., [1,1,1,1] }. The corresponding output thread_data in those threads will be { [1,2,3,4], [5,6,7,8], ..., [509,510,511,512] }. Furthermore the value 512 will be stored in block_aggregate for all threads.
Template Parameters
ITEMS_PER_THREAD[inferred] The number of consecutive items partitioned onto each thread.
ScanOp[inferred] Binary scan functor type having member T operator()(const T &a, const T &b)
Parameters
[in]inputCalling thread's input items
[out]outputCalling thread's output items (may be aliased to input)
[out]block_aggregateblock-wide aggregate reduction of input items
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
template<int ITEMS_PER_THREAD, typename BlockPrefixCallbackOp >
__device__ __forceinline__ void cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::InclusiveSum ( T(&)  input[ITEMS_PER_THREAD],
T(&)  output[ITEMS_PER_THREAD],
BlockPrefixCallbackOp &  block_prefix_callback_op 
)
inline

Computes an inclusive block-wide prefix scan using addition (+) as the scan operator. Each thread contributes an array of consecutive input elements. Instead of using 0 as the block-wide prefix, the call-back functor block_prefix_callback_op is invoked by the first warp in the block, and the value returned by lane0 in that warp is used as the "seed" value that logically prefixes the thread block's scan inputs. Also provides every thread with the block-wide block_aggregate of all inputs.

  • The block_prefix_callback_op functor must implement a member function T operator()(T block_aggregate). The functor's input parameter block_aggregate is the same value also returned by the scan operation. The functor will be invoked by the first warp of threads in the block, however only the return value from lane0 is applied as the block-wide prefix. Can be stateful.
  • Assumes a blocked arrangement of (block-threads*items-per-thread) items across the thread block, where threadi owns the ith 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.
Snippet
The code snippet below illustrates a single thread block that progressively computes an inclusive prefix sum over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 512 integer items that are partitioned in a blocked arrangement across 128 threads where each thread owns 4 consecutive items.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
// A stateful callback functor that maintains a running prefix to be applied
// during consecutive scan operations.
struct BlockPrefixCallbackOp
{
// Running prefix
int running_total;
// Constructor
__device__ BlockPrefixCallbackOp(int running_total) : running_total(running_total) {}
// Callback operator to be entered by the first warp of threads in the block.
// Thread-0 is responsible for returning a value for seeding the block-wide scan.
__device__ int operator()(int block_aggregate)
{
int old_prefix = running_total;
running_total += block_aggregate;
return old_prefix;
}
};
__global__ void ExampleKernel(int *d_data, int num_items, ...)
{
// Specialize BlockLoad, BlockStore, and BlockScan for a 1D block of 128 threads, 4 ints per thread
typedef cub::BlockLoad<int*, 128, 4, BLOCK_LOAD_TRANSPOSE> BlockLoad;
typedef cub::BlockStore<int, 128, 4, BLOCK_STORE_TRANSPOSE> BlockStore;
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate aliased shared memory for BlockLoad, BlockStore, and BlockScan
__shared__ union {
typename BlockLoad::TempStorage load;
typename BlockScan::TempStorage scan;
typename BlockStore::TempStorage store;
} temp_storage;
// Initialize running total
BlockPrefixCallbackOp prefix_op(0);
// Have the block iterate over segments of items
for (int block_offset = 0; block_offset < num_items; block_offset += 128 * 4)
{
// Load a segment of consecutive items that are blocked across threads
int thread_data[4];
BlockLoad(temp_storage.load).Load(d_data + block_offset, thread_data);
CTA_SYNC();
// Collectively compute the block-wide inclusive prefix sum
BlockScan(temp_storage.scan).IncluisveSum(
thread_data, thread_data, prefix_op);
CTA_SYNC();
// Store scanned items to output segment
BlockStore(temp_storage.store).Store(d_data + block_offset, thread_data);
CTA_SYNC();
}
Suppose the input d_data is 1, 1, 1, 1, 1, 1, 1, 1, .... The corresponding output for the first segment will be 1, 2, 3, 4, ..., 511, 512. The output for the second segment will be 513, 514, 515, 516, ..., 1023, 1024.
Template Parameters
ITEMS_PER_THREAD[inferred] The number of consecutive items partitioned onto each thread.
BlockPrefixCallbackOp[inferred] Call-back functor type having member T operator()(T block_aggregate)
Parameters
[in]inputCalling thread's input items
[out]outputCalling thread's output items (may be aliased to input)
[in,out]block_prefix_callback_op[warp0 only] Call-back functor for specifying a block-wide prefix to be applied to the logical input sequence.
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
template<typename ScanOp >
__device__ __forceinline__ void cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::InclusiveScan ( T  input,
T &  output,
ScanOp  scan_op 
)
inline

Computes an inclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes one input element.

  • Supports non-commutative scan operators.
  • For multi-dimensional blocks, threads are linearly ranked in row-major order.
  • 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.
Snippet
The code snippet below illustrates an inclusive prefix max scan of 128 integer items that are partitioned across 128 threads.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
__global__ void ExampleKernel(...)
{
// Specialize BlockScan for a 1D block of 128 threads of type int
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Obtain input item for each thread
int thread_data;
...
// Collectively compute the block-wide inclusive prefix max scan
BlockScan(temp_storage).InclusiveScan(thread_data, thread_data, cub::Max());
Suppose the set of input thread_data across the block of threads is 0, -1, 2, -3, ..., 126, -127. The corresponding output thread_data in those threads will be 0, 0, 2, 2, ..., 126, 126.
Template Parameters
ScanOp[inferred] Binary scan functor type having member T operator()(const T &a, const T &b)
Parameters
[in]inputCalling thread's input item
[out]outputCalling thread's output item (may be aliased to input)
[in]scan_opBinary scan functor
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
template<typename ScanOp >
__device__ __forceinline__ void cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::InclusiveScan ( T  input,
T &  output,
ScanOp  scan_op,
T &  block_aggregate 
)
inline

Computes an inclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes one input element. Also provides every thread with the block-wide block_aggregate of all inputs.

  • Supports non-commutative scan operators.
  • For multi-dimensional blocks, threads are linearly ranked in row-major order.
  • 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.
Snippet
The code snippet below illustrates an inclusive prefix max scan of 128 integer items that are partitioned across 128 threads.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
__global__ void ExampleKernel(...)
{
// Specialize BlockScan for a 1D block of 128 threads of type int
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Obtain input item for each thread
int thread_data;
...
// Collectively compute the block-wide inclusive prefix max scan
int block_aggregate;
BlockScan(temp_storage).InclusiveScan(thread_data, thread_data, cub::Max(), block_aggregate);
Suppose the set of input thread_data across the block of threads is 0, -1, 2, -3, ..., 126, -127. The corresponding output thread_data in those threads will be 0, 0, 2, 2, ..., 126, 126. Furthermore the value 126 will be stored in block_aggregate for all threads.
Template Parameters
ScanOp[inferred] Binary scan functor type having member T operator()(const T &a, const T &b)
Parameters
[in]inputCalling thread's input item
[out]outputCalling thread's output item (may be aliased to input)
[in]scan_opBinary scan functor
[out]block_aggregateblock-wide aggregate reduction of input items
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
template<typename ScanOp , typename BlockPrefixCallbackOp >
__device__ __forceinline__ void cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::InclusiveScan ( T  input,
T &  output,
ScanOp  scan_op,
BlockPrefixCallbackOp &  block_prefix_callback_op 
)
inline

Computes an inclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes one input element. the call-back functor block_prefix_callback_op is invoked by the first warp in the block, and the value returned by lane0 in that warp is used as the "seed" value that logically prefixes the thread block's scan inputs. Also provides every thread with the block-wide block_aggregate of all inputs.

  • The block_prefix_callback_op functor must implement a member function T operator()(T block_aggregate). The functor's input parameter block_aggregate is the same value also returned by the scan operation. The functor will be invoked by the first warp of threads in the block, however only the return value from lane0 is applied as the block-wide prefix. Can be stateful.
  • Supports non-commutative scan operators.
  • For multi-dimensional blocks, threads are linearly ranked in row-major order.
  • 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.
Snippet
The code snippet below illustrates a single thread block that progressively computes an inclusive prefix max scan over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 128 integer items that are partitioned across 128 threads.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
// A stateful callback functor that maintains a running prefix to be applied
// during consecutive scan operations.
struct BlockPrefixCallbackOp
{
// Running prefix
int running_total;
// Constructor
__device__ BlockPrefixCallbackOp(int running_total) : running_total(running_total) {}
// Callback operator to be entered by the first warp of threads in the block.
// Thread-0 is responsible for returning a value for seeding the block-wide scan.
__device__ int operator()(int block_aggregate)
{
int old_prefix = running_total;
running_total = (block_aggregate > old_prefix) ? block_aggregate : old_prefix;
return old_prefix;
}
};
__global__ void ExampleKernel(int *d_data, int num_items, ...)
{
// Specialize BlockScan for a 1D block of 128 threads
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Initialize running total
BlockPrefixCallbackOp prefix_op(INT_MIN);
// Have the block iterate over segments of items
for (int block_offset = 0; block_offset < num_items; block_offset += 128)
{
// Load a segment of consecutive items that are blocked across threads
int thread_data = d_data[block_offset];
// Collectively compute the block-wide inclusive prefix max scan
BlockScan(temp_storage).InclusiveScan(
thread_data, thread_data, cub::Max(), prefix_op);
CTA_SYNC();
// Store scanned items to output segment
d_data[block_offset] = thread_data;
}
Suppose the input d_data is 0, -1, 2, -3, 4, -5, .... The corresponding output for the first segment will be 0, 0, 2, 2, ..., 126, 126. The output for the second segment will be 128, 128, 130, 130, ..., 254, 254.
Template Parameters
ScanOp[inferred] Binary scan functor type having member T operator()(const T &a, const T &b)
BlockPrefixCallbackOp[inferred] Call-back functor type having member T operator()(T block_aggregate)
Parameters
[in]inputCalling thread's input item
[out]outputCalling thread's output item (may be aliased to input)
[in]scan_opBinary scan functor
[in,out]block_prefix_callback_op[warp0 only] Call-back functor for specifying a block-wide prefix to be applied to the logical input sequence.
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
template<int ITEMS_PER_THREAD, typename ScanOp >
__device__ __forceinline__ void cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::InclusiveScan ( T(&)  input[ITEMS_PER_THREAD],
T(&)  output[ITEMS_PER_THREAD],
ScanOp  scan_op 
)
inline

Computes an inclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes an array of consecutive input elements.

  • Supports non-commutative scan operators.
  • Assumes a blocked arrangement of (block-threads*items-per-thread) items across the thread block, where threadi owns the ith 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.
Snippet
The code snippet below illustrates an inclusive prefix max scan of 512 integer items that are partitioned in a blocked arrangement across 128 threads where each thread owns 4 consecutive items.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
__global__ void ExampleKernel(...)
{
// Specialize BlockScan for a 1D block of 128 threads of type int
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Obtain a segment of consecutive items that are blocked across threads
int thread_data[4];
...
// Collectively compute the block-wide inclusive prefix max scan
BlockScan(temp_storage).InclusiveScan(thread_data, thread_data, cub::Max());
Suppose the set of input thread_data across the block of threads is { [0,-1,2,-3], [4,-5,6,-7], ..., [508,-509,510,-511] }. The corresponding output thread_data in those threads will be { [0,0,2,2], [4,4,6,6], ..., [508,508,510,510] }.
Template Parameters
ITEMS_PER_THREAD[inferred] The number of consecutive items partitioned onto each thread.
ScanOp[inferred] Binary scan functor type having member T operator()(const T &a, const T &b)
Parameters
[in]inputCalling thread's input items
[out]outputCalling thread's output items (may be aliased to input)
[in]scan_opBinary scan functor
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
template<int ITEMS_PER_THREAD, typename ScanOp >
__device__ __forceinline__ void cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::InclusiveScan ( T(&)  input[ITEMS_PER_THREAD],
T(&)  output[ITEMS_PER_THREAD],
ScanOp  scan_op,
T &  block_aggregate 
)
inline

Computes an inclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes an array of consecutive input elements. Also provides every thread with the block-wide block_aggregate of all inputs.

  • Supports non-commutative scan operators.
  • Assumes a blocked arrangement of (block-threads*items-per-thread) items across the thread block, where threadi owns the ith 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.
Snippet
The code snippet below illustrates an inclusive prefix max scan of 512 integer items that are partitioned in a blocked arrangement across 128 threads where each thread owns 4 consecutive items.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
__global__ void ExampleKernel(...)
{
// Specialize BlockScan for a 1D block of 128 threads of type int
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate shared memory for BlockScan
__shared__ typename BlockScan::TempStorage temp_storage;
// Obtain a segment of consecutive items that are blocked across threads
int thread_data[4];
...
// Collectively compute the block-wide inclusive prefix max scan
int block_aggregate;
BlockScan(temp_storage).InclusiveScan(thread_data, thread_data, cub::Max(), block_aggregate);
Suppose the set of input thread_data across the block of threads is { [0,-1,2,-3], [4,-5,6,-7], ..., [508,-509,510,-511] }. The corresponding output thread_data in those threads will be { [0,0,2,2], [4,4,6,6], ..., [508,508,510,510] }. Furthermore the value 510 will be stored in block_aggregate for all threads.
Template Parameters
ITEMS_PER_THREAD[inferred] The number of consecutive items partitioned onto each thread.
ScanOp[inferred] Binary scan functor type having member T operator()(const T &a, const T &b)
Parameters
[in]inputCalling thread's input items
[out]outputCalling thread's output items (may be aliased to input)
[in]scan_opBinary scan functor
[out]block_aggregateblock-wide aggregate reduction of input items
template<typename T , int BLOCK_DIM_X, BlockScanAlgorithm ALGORITHM = BLOCK_SCAN_RAKING, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
template<int ITEMS_PER_THREAD, typename ScanOp , typename BlockPrefixCallbackOp >
__device__ __forceinline__ void cub::BlockScan< T, BLOCK_DIM_X, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::InclusiveScan ( T(&)  input[ITEMS_PER_THREAD],
T(&)  output[ITEMS_PER_THREAD],
ScanOp  scan_op,
BlockPrefixCallbackOp &  block_prefix_callback_op 
)
inline

Computes an inclusive block-wide prefix scan using the specified binary scan_op functor. Each thread contributes an array of consecutive input elements. the call-back functor block_prefix_callback_op is invoked by the first warp in the block, and the value returned by lane0 in that warp is used as the "seed" value that logically prefixes the thread block's scan inputs. Also provides every thread with the block-wide block_aggregate of all inputs.

  • The block_prefix_callback_op functor must implement a member function T operator()(T block_aggregate). The functor's input parameter block_aggregate is the same value also returned by the scan operation. The functor will be invoked by the first warp of threads in the block, however only the return value from lane0 is applied as the block-wide prefix. Can be stateful.
  • Supports non-commutative scan operators.
  • Assumes a blocked arrangement of (block-threads*items-per-thread) items across the thread block, where threadi owns the ith 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.
Snippet
The code snippet below illustrates a single thread block that progressively computes an inclusive prefix max scan over multiple "tiles" of input using a prefix functor to maintain a running total between block-wide scans. Each tile consists of 128 integer items that are partitioned across 128 threads.
#include <cub/cub.cuh> // or equivalently <cub/block/block_scan.cuh>
// A stateful callback functor that maintains a running prefix to be applied
// during consecutive scan operations.
struct BlockPrefixCallbackOp
{
// Running prefix
int running_total;
// Constructor
__device__ BlockPrefixCallbackOp(int running_total) : running_total(running_total) {}
// Callback operator to be entered by the first warp of threads in the block.
// Thread-0 is responsible for returning a value for seeding the block-wide scan.
__device__ int operator()(int block_aggregate)
{
int old_prefix = running_total;
running_total = (block_aggregate > old_prefix) ? block_aggregate : old_prefix;
return old_prefix;
}
};
__global__ void ExampleKernel(int *d_data, int num_items, ...)
{
// Specialize BlockLoad, BlockStore, and BlockScan for a 1D block of 128 threads, 4 ints per thread
typedef cub::BlockLoad<int*, 128, 4, BLOCK_LOAD_TRANSPOSE> BlockLoad;
typedef cub::BlockStore<int, 128, 4, BLOCK_STORE_TRANSPOSE> BlockStore;
typedef cub::BlockScan<int, 128> BlockScan;
// Allocate aliased shared memory for BlockLoad, BlockStore, and BlockScan
__shared__ union {
typename BlockLoad::TempStorage load;
typename BlockScan::TempStorage scan;
typename BlockStore::TempStorage store;
} temp_storage;
// Initialize running total
BlockPrefixCallbackOp prefix_op(0);
// Have the block iterate over segments of items
for (int block_offset = 0; block_offset < num_items; block_offset += 128 * 4)
{
// Load a segment of consecutive items that are blocked across threads
int thread_data[4];
BlockLoad(temp_storage.load).Load(d_data + block_offset, thread_data);
CTA_SYNC();
// Collectively compute the block-wide inclusive prefix max scan
BlockScan(temp_storage.scan).InclusiveScan(
thread_data, thread_data, cub::Max(), prefix_op);
CTA_SYNC();
// Store scanned items to output segment
BlockStore(temp_storage.store).Store(d_data + block_offset, thread_data);
CTA_SYNC();
}
Suppose the input d_data is 0, -1, 2, -3, 4, -5, .... The corresponding output for the first segment will be 0, 0, 2, 2, 4, 4, ..., 510, 510. The output for the second segment will be 512, 512, 514, 514, 516, 516, ..., 1022, 1022.
Template Parameters
ITEMS_PER_THREAD[inferred] The number of consecutive items partitioned onto each thread.
ScanOp[inferred] Binary scan functor type having member T operator()(const T &a, const T &b)
BlockPrefixCallbackOp[inferred] Call-back functor type having member T operator()(T block_aggregate)
Parameters
[in]inputCalling thread's input items
[out]outputCalling thread's output items (may be aliased to input)
[in]scan_opBinary scan functor
[in,out]block_prefix_callback_op[warp0 only] Call-back functor for specifying a block-wide prefix to be applied to the logical input sequence.
The documentation for this class was generated from the following file:
Generated by   doxygen 1.8.4
© 2022 NVIDIA Corporation