CUDPP  2.1
CUDA Data-Parallel Primitives Library
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Groups Pages
Public Types | Static Public Member Functions | List of all members
SegmentedScanTraits< T, Oper, backward, exclusivity, doShiftFlags, fullBlock, sums, sm12OrBetter > Class Template Reference

Template class containing compile-time parameters to the segmented scan functions. More...

Public Types

typedef Oper Op
 The operator functor used for segmented scan.
 

Static Public Member Functions

static __device__ bool isBackward ()
 
static __device__ bool isExclusive ()
 
static __device__ bool shiftFlags ()
 
static __device__ bool isFullBlock ()
 
static __device__ bool writeSums ()
 
static __device__ bool isSM12OrBetter ()
 

Detailed Description

template<typename T, class Oper, bool backward, bool exclusivity, bool doShiftFlags, bool fullBlock, bool sums, bool sm12OrBetter>
class SegmentedScanTraits< T, Oper, backward, exclusivity, doShiftFlags, fullBlock, sums, sm12OrBetter >

Template class containing compile-time parameters to the segmented scan functions.

SegmentedScanTraits is passed as a template parameter to all segmented scan functions. By using these compile-time functions we can enable generic code while maintaining the highest performance. This is crucial for the performance of low-level workhorse algorithms like segmented scan.

Parameters
TThe datatype of the segmented scan
operThe CUDPPOperator to use for the segmented scan (add, max, etc.)
unrollTrue if scan inner loops should be unrolled
sumsTrue if each block should write it's sum to the d_blockSums array (false for single-block scans)
backwardTrue if this is a backward scan, False if this is a forward scan
fullBlockTrue if all blocks in this scan are full (CTA_SIZE * SCAN_ELEMENTS_PER_THREAD elements)
exclusivityTrue for exclusive scans, false for inclusive scans

Member Function Documentation

template<typename T , class Oper , bool backward, bool exclusivity, bool doShiftFlags, bool fullBlock, bool sums, bool sm12OrBetter>
static __device__ bool SegmentedScanTraits< T, Oper, backward, exclusivity, doShiftFlags, fullBlock, sums, sm12OrBetter >::isBackward ( )
inlinestatic
Returns
true if this is a backward scan
template<typename T , class Oper , bool backward, bool exclusivity, bool doShiftFlags, bool fullBlock, bool sums, bool sm12OrBetter>
static __device__ bool SegmentedScanTraits< T, Oper, backward, exclusivity, doShiftFlags, fullBlock, sums, sm12OrBetter >::isExclusive ( )
inlinestatic
Returns
true if this is an exclusive scan
template<typename T , class Oper , bool backward, bool exclusivity, bool doShiftFlags, bool fullBlock, bool sums, bool sm12OrBetter>
static __device__ bool SegmentedScanTraits< T, Oper, backward, exclusivity, doShiftFlags, fullBlock, sums, sm12OrBetter >::shiftFlags ( )
inlinestatic
Returns
true if this scan needs to shift flags to the left. This is only needed for the first level scan in a multi-block scan
template<typename T , class Oper , bool backward, bool exclusivity, bool doShiftFlags, bool fullBlock, bool sums, bool sm12OrBetter>
static __device__ bool SegmentedScanTraits< T, Oper, backward, exclusivity, doShiftFlags, fullBlock, sums, sm12OrBetter >::isFullBlock ( )
inlinestatic
Returns
true if this is a full scan – all blocks process CTA_SIZE * SCAN_ELEMENTS_PER_THREAD elements
template<typename T , class Oper , bool backward, bool exclusivity, bool doShiftFlags, bool fullBlock, bool sums, bool sm12OrBetter>
static __device__ bool SegmentedScanTraits< T, Oper, backward, exclusivity, doShiftFlags, fullBlock, sums, sm12OrBetter >::writeSums ( )
inlinestatic
Returns
true if this scan writes the sum of each block to the d_blockSums array (multi-block scans)
template<typename T , class Oper , bool backward, bool exclusivity, bool doShiftFlags, bool fullBlock, bool sums, bool sm12OrBetter>
static __device__ bool SegmentedScanTraits< T, Oper, backward, exclusivity, doShiftFlags, fullBlock, sums, sm12OrBetter >::isSM12OrBetter ( )
inlinestatic
Returns
true if we are sm12 or better hardware

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