CUB
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Groups
List of all members
cub::DeviceScan Struct Reference

Detailed description

DeviceScan provides device-wide, parallel operations for computing a prefix scan across a sequence of data items residing within device-accessible memory.

device_scan.png
.
Overview
Given a sequence of input elements and a binary reduction operator, a prefix scan produces an output sequence where each element is computed to be the reduction of the elements occurring earlier in the input sequence. 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.
Usage Considerations
  • Dynamic parallelism. DeviceScan methods can be called within kernel code on devices in which CUDA dynamic parallelism is supported.
Performance
The work-complexity of prefix scan as a function of input size is linear, resulting in performance throughput that plateaus with problem sizes large enough to saturate the GPU.
The following chart illustrates DeviceScan::ExclusiveSum performance across different CUDA architectures for int32 keys. Performance plots for other scenarios can be found in the detailed method descriptions below.
scan_int32.png

Definition at line 77 of file device_scan.cuh.

Static Public Methods

Exclusive scans
template<typename InputIteratorT , typename OutputIteratorT >
static CUB_RUNTIME_FUNCTION
cudaError_t 
ExclusiveSum (void *d_temp_storage, size_t &temp_storage_bytes, InputIteratorT d_in, OutputIteratorT d_out, int num_items, cudaStream_t stream=0, bool debug_synchronous=false)
 Computes a device-wide exclusive prefix sum. More...
 
template<typename InputIteratorT , typename OutputIteratorT , typename ScanOp , typename Identity >
static CUB_RUNTIME_FUNCTION
cudaError_t 
ExclusiveScan (void *d_temp_storage, size_t &temp_storage_bytes, InputIteratorT d_in, OutputIteratorT d_out, ScanOp scan_op, Identity identity, int num_items, cudaStream_t stream=0, bool debug_synchronous=false)
 Computes a device-wide exclusive prefix scan using the specified binary scan_op functor. More...
 
Inclusive scans
template<typename InputIteratorT , typename OutputIteratorT >
static CUB_RUNTIME_FUNCTION
cudaError_t 
InclusiveSum (void *d_temp_storage, size_t &temp_storage_bytes, InputIteratorT d_in, OutputIteratorT d_out, int num_items, cudaStream_t stream=0, bool debug_synchronous=false)
 Computes a device-wide inclusive prefix sum. More...
 
template<typename InputIteratorT , typename OutputIteratorT , typename ScanOp >
static CUB_RUNTIME_FUNCTION
cudaError_t 
InclusiveScan (void *d_temp_storage, size_t &temp_storage_bytes, InputIteratorT d_in, OutputIteratorT d_out, ScanOp scan_op, int num_items, cudaStream_t stream=0, bool debug_synchronous=false)
 Computes a device-wide inclusive prefix scan using the specified binary scan_op functor. More...
 

Member Function Documentation

template<typename InputIteratorT , typename OutputIteratorT >
static CUB_RUNTIME_FUNCTION cudaError_t cub::DeviceScan::ExclusiveSum ( void *  d_temp_storage,
size_t &  temp_storage_bytes,
InputIteratorT  d_in,
OutputIteratorT  d_out,
int  num_items,
cudaStream_t  stream = 0,
bool  debug_synchronous = false 
)
inlinestatic

Computes a device-wide exclusive prefix sum.

  • Supports non-commutative sum operators.
  • When d_temp_storage is NULL, no work is done and the required allocation size is returned in temp_storage_bytes.
Performance
The following charts illustrate saturated exclusive sum performance across different CUDA architectures for int32 and int64 items, respectively.
scan_int32.png
scan_int64.png
Snippet
The code snippet below illustrates the exclusive prefix sum of an int device vector.
#include <cub/cub.cuh> // or equivalently <cub/device/device_scan.cuh>
// Declare, allocate, and initialize device-accessible pointers for input and output
int num_items; // e.g., 7
int *d_in; // e.g., [8, 6, 7, 5, 3, 0, 9]
int *d_out; // e.g., [ , , , , , , ]
...
// Determine temporary device storage requirements
void *d_temp_storage = NULL;
size_t temp_storage_bytes = 0;
cub::DeviceScan::ExclusiveSum(d_temp_storage, temp_storage_bytes, d_in, d_out, num_items);
// Allocate temporary storage
cudaMalloc(&d_temp_storage, temp_storage_bytes);
// Run exclusive prefix sum
cub::DeviceScan::ExclusiveSum(d_temp_storage, temp_storage_bytes, d_in, d_out, num_items);
// d_out s<-- [0, 8, 14, 21, 26, 29, 29]
Template Parameters
InputIteratorT[inferred] Random-access input iterator type for reading scan inputs (may be a simple pointer type)
OutputIteratorT[inferred] Random-access output iterator type for writing scan outputs (may be a simple pointer type)
Parameters
[in]d_temp_storageDevice-accessible allocation of temporary storage. When NULL, the required allocation size is written to temp_storage_bytes and no work is done.
[in,out]temp_storage_bytesReference to size in bytes of d_temp_storage allocation
[in]d_inPointer to the input sequence of data items
[out]d_outPointer to the output sequence of data items
[in]num_itemsTotal number of input items (i.e., the length of d_in)
[in]stream[optional] CUDA stream to launch kernels within. Default is stream0.
[in]debug_synchronous[optional] Whether or not to synchronize the stream after every kernel launch to check for errors. May cause significant slowdown. Default is false.
Examples:
example_device_scan.cu.

Definition at line 132 of file device_scan.cuh.

template<typename InputIteratorT , typename OutputIteratorT , typename ScanOp , typename Identity >
static CUB_RUNTIME_FUNCTION cudaError_t cub::DeviceScan::ExclusiveScan ( void *  d_temp_storage,
size_t &  temp_storage_bytes,
InputIteratorT  d_in,
OutputIteratorT  d_out,
ScanOp  scan_op,
Identity  identity,
int  num_items,
cudaStream_t  stream = 0,
bool  debug_synchronous = false 
)
inlinestatic

Computes a device-wide exclusive prefix scan using the specified binary scan_op functor.

  • Supports non-commutative scan operators.
  • When d_temp_storage is NULL, no work is done and the required allocation size is returned in temp_storage_bytes.
Snippet
The code snippet below illustrates the exclusive prefix min-scan of an int device vector
#include <cub/cub.cuh> // or equivalently <cub/device/device_scan.cuh>
// CustomMin functor
struct CustomMin
{
template <typename T>
CUB_RUNTIME_FUNCTION __forceinline__
T operator()(const T &a, const T &b) const {
return (b < a) ? b : a;
}
};
// Declare, allocate, and initialize device-accessible pointers for input and output
int num_items; // e.g., 7
int *d_in; // e.g., [8, 6, 7, 5, 3, 0, 9]
int *d_out; // e.g., [ , , , , , , ]
CustomMin min_op
...
// Determine temporary device storage requirements for exclusive prefix scan
void *d_temp_storage = NULL;
size_t temp_storage_bytes = 0;
cub::DeviceScan::ExclusiveScan(d_temp_storage, temp_storage_bytes, d_in, d_out, min_op, (int) MAX_INT, num_items);
// Allocate temporary storage for exclusive prefix scan
cudaMalloc(&d_temp_storage, temp_storage_bytes);
// Run exclusive prefix min-scan
cub::DeviceScan::ExclusiveScan(d_temp_storage, temp_storage_bytes, d_in, d_out, min_op, (int) MAX_INT, num_items);
// d_out <-- [2147483647, 8, 6, 6, 5, 3, 0]
Template Parameters
InputIteratorT[inferred] Random-access input iterator type for reading scan inputs (may be a simple pointer type)
OutputIteratorT[inferred] Random-access output iterator type for writing scan outputs (may be a simple pointer type)
ScanOp[inferred] Binary scan functor type having member T operator()(const T &a, const T &b)
Identity[inferred] Type of the identity value used Binary scan functor type having member T operator()(const T &a, const T &b)
Parameters
[in]d_temp_storageDevice-accessible allocation of temporary storage. When NULL, the required allocation size is written to temp_storage_bytes and no work is done.
[in,out]temp_storage_bytesReference to size in bytes of d_temp_storage allocation
[in]d_inPointer to the input sequence of data items
[out]d_outPointer to the output sequence of data items
[in]scan_opBinary scan functor
[in]identityIdentity element
[in]num_itemsTotal number of input items (i.e., the length of d_in)
[in]stream[optional] CUDA stream to launch kernels within. Default is stream0.
[in]debug_synchronous[optional] Whether or not to synchronize the stream after every kernel launch to check for errors. May cause significant slowdown. Default is false.

Definition at line 216 of file device_scan.cuh.

template<typename InputIteratorT , typename OutputIteratorT >
static CUB_RUNTIME_FUNCTION cudaError_t cub::DeviceScan::InclusiveSum ( void *  d_temp_storage,
size_t &  temp_storage_bytes,
InputIteratorT  d_in,
OutputIteratorT  d_out,
int  num_items,
cudaStream_t  stream = 0,
bool  debug_synchronous = false 
)
inlinestatic

Computes a device-wide inclusive prefix sum.

  • Supports non-commutative sum operators.
  • When d_temp_storage is NULL, no work is done and the required allocation size is returned in temp_storage_bytes.
Snippet
The code snippet below illustrates the inclusive prefix sum of an int device vector.
#include <cub/cub.cuh> // or equivalently <cub/device/device_scan.cuh>
// Declare, allocate, and initialize device-accessible pointers for input and output
int num_items; // e.g., 7
int *d_in; // e.g., [8, 6, 7, 5, 3, 0, 9]
int *d_out; // e.g., [ , , , , , , ]
...
// Determine temporary device storage requirements for inclusive prefix sum
void *d_temp_storage = NULL;
size_t temp_storage_bytes = 0;
cub::DeviceScan::InclusiveSum(d_temp_storage, temp_storage_bytes, d_in, d_out, num_items);
// Allocate temporary storage for inclusive prefix sum
cudaMalloc(&d_temp_storage, temp_storage_bytes);
// Run inclusive prefix sum
cub::DeviceScan::InclusiveSum(d_temp_storage, temp_storage_bytes, d_in, d_out, num_items);
// d_out <-- [8, 14, 21, 26, 29, 29, 38]
Template Parameters
InputIteratorT[inferred] Random-access input iterator type for reading scan inputs (may be a simple pointer type)
OutputIteratorT[inferred] Random-access output iterator type for writing scan outputs (may be a simple pointer type)
Parameters
[in]d_temp_storageDevice-accessible allocation of temporary storage. When NULL, the required allocation size is written to temp_storage_bytes and no work is done.
[in,out]temp_storage_bytesReference to size in bytes of d_temp_storage allocation
[in]d_inPointer to the input sequence of data items
[out]d_outPointer to the output sequence of data items
[in]num_itemsTotal number of input items (i.e., the length of d_in)
[in]stream[optional] CUDA stream to launch kernels within. Default is stream0.
[in]debug_synchronous[optional] Whether or not to synchronize the stream after every kernel launch to check for errors. May cause significant slowdown. Default is false.

Definition at line 291 of file device_scan.cuh.

template<typename InputIteratorT , typename OutputIteratorT , typename ScanOp >
static CUB_RUNTIME_FUNCTION cudaError_t cub::DeviceScan::InclusiveScan ( void *  d_temp_storage,
size_t &  temp_storage_bytes,
InputIteratorT  d_in,
OutputIteratorT  d_out,
ScanOp  scan_op,
int  num_items,
cudaStream_t  stream = 0,
bool  debug_synchronous = false 
)
inlinestatic

Computes a device-wide inclusive prefix scan using the specified binary scan_op functor.

  • Supports non-commutative scan operators.
  • When d_temp_storage is NULL, no work is done and the required allocation size is returned in temp_storage_bytes.
Snippet
The code snippet below illustrates the inclusive prefix min-scan of an int device vector.
#include <cub/cub.cuh> // or equivalently <cub/device/device_scan.cuh>
// CustomMin functor
struct CustomMin
{
template <typename T>
CUB_RUNTIME_FUNCTION __forceinline__
T operator()(const T &a, const T &b) const {
return (b < a) ? b : a;
}
};
// Declare, allocate, and initialize device-accessible pointers for input and output
int num_items; // e.g., 7
int *d_in; // e.g., [8, 6, 7, 5, 3, 0, 9]
int *d_out; // e.g., [ , , , , , , ]
CustomMin min_op;
...
// Determine temporary device storage requirements for inclusive prefix scan
void *d_temp_storage = NULL;
size_t temp_storage_bytes = 0;
cub::DeviceScan::InclusiveScan(d_temp_storage, temp_storage_bytes, d_in, d_out, min_op, num_items);
// Allocate temporary storage for inclusive prefix scan
cudaMalloc(&d_temp_storage, temp_storage_bytes);
// Run inclusive prefix min-scan
cub::DeviceScan::InclusiveScan(d_temp_storage, temp_storage_bytes, d_in, d_out, min_op, num_items);
// d_out <-- [8, 6, 6, 5, 3, 0, 0]
Template Parameters
InputIteratorT[inferred] Random-access input iterator type for reading scan inputs (may be a simple pointer type)
OutputIteratorT[inferred] Random-access output iterator type for writing scan outputs (may be a simple pointer type)
ScanOp[inferred] Binary scan functor type having member T operator()(const T &a, const T &b)
Parameters
[in]d_temp_storageDevice-accessible allocation of temporary storage. When NULL, the required allocation size is written to temp_storage_bytes and no work is done.
[in,out]temp_storage_bytesReference to size in bytes of d_temp_storage allocation
[in]d_inPointer to the input sequence of data items
[out]d_outPointer to the output sequence of data items
[in]scan_opBinary scan functor
[in]num_itemsTotal number of input items (i.e., the length of d_in)
[in]stream[optional] CUDA stream to launch kernels within. Default is stream0.
[in]debug_synchronous[optional] Whether or not to synchronize the stream after every kernel launch to check for errors. May cause significant slowdown. Default is false.

Definition at line 370 of file device_scan.cuh.


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