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.
|
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, ScanTileStateT > | TilePrefixCallbackOpT |
|
typedef BlockScanRunningPrefixOp< OutputT, ScanOpT > | RunningPrefixCallbackOp |
|
|
__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) |
|
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_items | Total number of input items |
tile_state | Global tile state descriptor |
start_tile | The starting tile for the current grid |
Definition at line 333 of file agent_scan.cuh.
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_offset | Threadblock begin offset (inclusive) |
[in] | range_end | Threadblock end offset (exclusive) |
Definition at line 406 of file agent_scan.cuh.
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_offset | Threadblock begin offset (inclusive) |
[in] | range_end | Threadblock end offset (exclusive) |
[in] | prefix | The prefix to apply to the scan segment |
Definition at line 444 of file agent_scan.cuh.
template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >
template<bool IS_LAST_TILE>
< Whether the current tile is the last tile
Process a tile of input (dynamic chained scan)
- Parameters
-
num_remaining | Number of global input items remaining (including this tile) |
tile_idx | Tile index |
tile_offset | Tile offset |
tile_state | Global tile state descriptor |
Definition at line 288 of file agent_scan.cuh.
template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >
template<bool IS_FIRST_TILE, bool IS_LAST_TILE>
Process a tile of input
- Parameters
-
tile_offset | Tile offset |
prefix_op | Running prefix operator |
valid_items | Number of valid items in the tile |
Definition at line 366 of file agent_scan.cuh.
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.
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.
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.
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.