OpenFPM_pdata  4.1.0
Project that contain the implementation of distributed structures
cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT > Struct Template Reference

AgentScan implements a stateful abstraction of CUDA thread blocks for participating in device-wide prefix scan . More...

Detailed Description

template<typename AgentScanPolicyT, typename InputIteratorT, typename OutputIteratorT, typename ScanOpT, typename InitValueT, typename OffsetT>
struct cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >

AgentScan implements a stateful abstraction of CUDA thread blocks for participating in device-wide prefix scan .

< Signed integer type for global offsets

Definition at line 98 of file agent_scan.cuh.

Data Structures

union  _TempStorage
 
struct  TempStorage
 

Public Types

enum  { IS_INCLUSIVE = Equals<InitValueT, NullType>::VALUE, BLOCK_THREADS = AgentScanPolicyT::BLOCK_THREADS, ITEMS_PER_THREAD = AgentScanPolicyT::ITEMS_PER_THREAD, TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD }
 
typedef std::iterator_traits< InputIteratorT >::value_type InputT
 
typedef If<(Equals< typename std::iterator_traits< OutputIteratorT >::value_type, void >::VALUE), typename std::iterator_traits< InputIteratorT >::value_type, typename std::iterator_traits< OutputIteratorT >::value_type >::Type OutputT
 
typedef ScanTileState< OutputT > ScanTileStateT
 
typedef If< IsPointer< InputIteratorT >::VALUE, CacheModifiedInputIterator< AgentScanPolicyT::LOAD_MODIFIER, InputT, OffsetT >, InputIteratorT >::Type WrappedInputIteratorT
 
typedef BlockLoad< OutputT, AgentScanPolicyT::BLOCK_THREADS, AgentScanPolicyT::ITEMS_PER_THREAD, AgentScanPolicyT::LOAD_ALGORITHM > BlockLoadT
 
typedef BlockStore< OutputT, AgentScanPolicyT::BLOCK_THREADS, AgentScanPolicyT::ITEMS_PER_THREAD, AgentScanPolicyT::STORE_ALGORITHM > BlockStoreT
 
typedef BlockScan< OutputT, AgentScanPolicyT::BLOCK_THREADS, AgentScanPolicyT::SCAN_ALGORITHM > BlockScanT
 
typedef TilePrefixCallbackOp< OutputT, ScanOpT, ScanTileStateTTilePrefixCallbackOpT
 
typedef BlockScanRunningPrefixOp< OutputT, ScanOpT > RunningPrefixCallbackOp
 

Public Member Functions

__device__ __forceinline__ void ScanTile (OutputT(&items)[ITEMS_PER_THREAD], OutputT init_value, ScanOpT scan_op, OutputT &block_aggregate, Int2Type< false >)
 
__device__ __forceinline__ void ScanTile (OutputT(&items)[ITEMS_PER_THREAD], InitValueT, ScanOpT scan_op, OutputT &block_aggregate, Int2Type< true >)
 
template<typename PrefixCallback >
__device__ __forceinline__ void ScanTile (OutputT(&items)[ITEMS_PER_THREAD], ScanOpT scan_op, PrefixCallback &prefix_op, Int2Type< false >)
 
template<typename PrefixCallback >
__device__ __forceinline__ void ScanTile (OutputT(&items)[ITEMS_PER_THREAD], ScanOpT scan_op, PrefixCallback &prefix_op, Int2Type< true >)
 
__device__ __forceinline__ AgentScan (TempStorage &temp_storage, InputIteratorT d_in, OutputIteratorT d_out, ScanOpT scan_op, InitValueT init_value)
 
template<bool IS_LAST_TILE>
__device__ __forceinline__ void ConsumeTile (OffsetT num_remaining, int tile_idx, OffsetT tile_offset, ScanTileStateT &tile_state)
 < Whether the current tile is the last tile More...
 
__device__ __forceinline__ void ConsumeRange (int num_items, ScanTileStateT &tile_state, int start_tile)
 
template<bool IS_FIRST_TILE, bool IS_LAST_TILE>
__device__ __forceinline__ void ConsumeTile (OffsetT tile_offset, RunningPrefixCallbackOp &prefix_op, int valid_items=TILE_ITEMS)
 
__device__ __forceinline__ void ConsumeRange (OffsetT range_offset, OffsetT range_end)
 
__device__ __forceinline__ void ConsumeRange (OffsetT range_offset, OffsetT range_end, OutputT prefix)
 

Data Fields

_TempStoragetemp_storage
 Reference to temp_storage.
 
WrappedInputIteratorT d_in
 Input data.
 
OutputIteratorT d_out
 Output data.
 
ScanOpT scan_op
 Binary scan operator.
 
InitValueT init_value
 The init_value element for ScanOpT.
 

Constructor & Destructor Documentation

◆ AgentScan()

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >
__device__ __forceinline__ cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::AgentScan ( TempStorage temp_storage,
InputIteratorT  d_in,
OutputIteratorT  d_out,
ScanOpT  scan_op,
InitValueT  init_value 
)
inline
Parameters
temp_storageReference to temp_storage
d_inInput data
d_outOutput data
scan_opBinary scan operator
init_valueInitial value to seed the exclusive scan

Definition at line 265 of file agent_scan.cuh.

Member Function Documentation

◆ ConsumeRange() [1/3]

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >
__device__ __forceinline__ void cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::ConsumeRange ( int  num_items,
ScanTileStateT tile_state,
int  start_tile 
)
inline

Scan tiles of items as part of a dynamic chained scan

Parameters
num_itemsTotal number of input items
tile_stateGlobal tile state descriptor
start_tileThe starting tile for the current grid

Definition at line 333 of file agent_scan.cuh.

◆ ConsumeRange() [2/3]

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >
__device__ __forceinline__ void cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::ConsumeRange ( OffsetT  range_offset,
OffsetT  range_end 
)
inline

Scan a consecutive share of input tiles

Parameters
[in]range_offsetThreadblock begin offset (inclusive)
[in]range_endThreadblock end offset (exclusive)

Definition at line 406 of file agent_scan.cuh.

◆ ConsumeRange() [3/3]

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >
__device__ __forceinline__ void cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::ConsumeRange ( OffsetT  range_offset,
OffsetT  range_end,
OutputT  prefix 
)
inline

Scan a consecutive share of input tiles, seeded with the specified prefix value

Parameters
[in]range_offsetThreadblock begin offset (inclusive)
[in]range_endThreadblock end offset (exclusive)
[in]prefixThe prefix to apply to the scan segment

Definition at line 444 of file agent_scan.cuh.

◆ ConsumeTile() [1/2]

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >
template<bool IS_LAST_TILE>
__device__ __forceinline__ void cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::ConsumeTile ( OffsetT  num_remaining,
int  tile_idx,
OffsetT  tile_offset,
ScanTileStateT tile_state 
)
inline

< Whether the current tile is the last tile

Process a tile of input (dynamic chained scan)

Parameters
num_remainingNumber of global input items remaining (including this tile)
tile_idxTile index
tile_offsetTile offset
tile_stateGlobal tile state descriptor

Definition at line 288 of file agent_scan.cuh.

◆ ConsumeTile() [2/2]

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >
template<bool IS_FIRST_TILE, bool IS_LAST_TILE>
__device__ __forceinline__ void cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::ConsumeTile ( OffsetT  tile_offset,
RunningPrefixCallbackOp prefix_op,
int  valid_items = TILE_ITEMS 
)
inline

Process a tile of input

Parameters
tile_offsetTile offset
prefix_opRunning prefix operator
valid_itemsNumber of valid items in the tile

Definition at line 366 of file agent_scan.cuh.

◆ ScanTile() [1/4]

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >
__device__ __forceinline__ void cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::ScanTile ( OutputT(&)  items[ITEMS_PER_THREAD],
OutputT  init_value,
ScanOpT  scan_op,
OutputT &  block_aggregate,
Int2Type< false >   
)
inline

Exclusive scan specialization (first tile)

Definition at line 202 of file agent_scan.cuh.

◆ ScanTile() [2/4]

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >
__device__ __forceinline__ void cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::ScanTile ( OutputT(&)  items[ITEMS_PER_THREAD],
InitValueT  ,
ScanOpT  scan_op,
OutputT &  block_aggregate,
Int2Type< true >   
)
inline

Inclusive scan specialization (first tile)

Definition at line 218 of file agent_scan.cuh.

◆ ScanTile() [3/4]

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >
template<typename PrefixCallback >
__device__ __forceinline__ void cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::ScanTile ( OutputT(&)  items[ITEMS_PER_THREAD],
ScanOpT  scan_op,
PrefixCallback &  prefix_op,
Int2Type< false >   
)
inline

Exclusive scan specialization (subsequent tiles)

Definition at line 234 of file agent_scan.cuh.

◆ ScanTile() [4/4]

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >
template<typename PrefixCallback >
__device__ __forceinline__ void cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::ScanTile ( OutputT(&)  items[ITEMS_PER_THREAD],
ScanOpT  scan_op,
PrefixCallback &  prefix_op,
Int2Type< true >   
)
inline

Inclusive scan specialization (subsequent tiles)

Definition at line 249 of file agent_scan.cuh.


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