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< typenamestd::iterator_traits< OutputIteratorT >::value_type, void >::VALUE), typenamestd::iterator_traits< InputIteratorT >::value_type, typenamestd::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

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

__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
_TempStorage &	temp_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.

Member Typedef Documentation

◆ BlockLoadT

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >

typedef BlockLoad< OutputT, AgentScanPolicyT::BLOCK_THREADS, AgentScanPolicyT::ITEMS_PER_THREAD, AgentScanPolicyT::LOAD_ALGORITHM> cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::BlockLoadT

Definition at line 136 of file agent_scan.cuh.

◆ BlockScanT

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >

typedef BlockScan< OutputT, AgentScanPolicyT::BLOCK_THREADS, AgentScanPolicyT::SCAN_ALGORITHM> cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::BlockScanT

Definition at line 151 of file agent_scan.cuh.

◆ BlockStoreT

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >

typedef BlockStore< OutputT, AgentScanPolicyT::BLOCK_THREADS, AgentScanPolicyT::ITEMS_PER_THREAD, AgentScanPolicyT::STORE_ALGORITHM> cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::BlockStoreT

Definition at line 144 of file agent_scan.cuh.

◆ InputT

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >

typedef std::iterator_traits<InputIteratorT>::value_type cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::InputT

Definition at line 105 of file agent_scan.cuh.

◆ OutputT

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >

typedef If<(Equals<typenamestd::iterator_traits<OutputIteratorT>::value_type,void>::VALUE),typenamestd::iterator_traits<InputIteratorT>::value_type,typenamestd::iterator_traits<OutputIteratorT>::value_type>::Type cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::OutputT

Definition at line 110 of file agent_scan.cuh.

◆ RunningPrefixCallbackOp

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >

typedef BlockScanRunningPrefixOp< OutputT, ScanOpT> cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::RunningPrefixCallbackOp

Definition at line 164 of file agent_scan.cuh.

◆ ScanTileStateT

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >

typedef ScanTileState<OutputT> cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::ScanTileStateT

Definition at line 113 of file agent_scan.cuh.

◆ TilePrefixCallbackOpT

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >

typedef TilePrefixCallbackOp< OutputT, ScanOpT, ScanTileStateT> cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::TilePrefixCallbackOpT

Definition at line 158 of file agent_scan.cuh.

◆ WrappedInputIteratorT

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >

typedef If<IsPointer<InputIteratorT>::VALUE,CacheModifiedInputIterator<AgentScanPolicyT::LOAD_MODIFIER,InputT,OffsetT>,InputIteratorT>::Type cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::WrappedInputIteratorT

Definition at line 119 of file agent_scan.cuh.

Member Enumeration Documentation

◆ anonymous enum

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >

anonymous enum

Definition at line 122 of file agent_scan.cuh.

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_storage	Reference to temp_storage
d_in	Input data
d_out	Output data
scan_op	Binary scan operator
init_value	Initial 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_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.

◆ 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_offset	Threadblock begin offset (inclusive)
[in]	range_end	Threadblock 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_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.

◆ 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_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.

◆ 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_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.

◆ 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],
		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() [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],
		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() [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.

Field Documentation

◆ d_in

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >

WrappedInputIteratorT cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::d_in

Input data.

Definition at line 188 of file agent_scan.cuh.

◆ d_out

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >

OutputIteratorT cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::d_out

Output data.

Definition at line 189 of file agent_scan.cuh.

◆ init_value

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >

InitValueT cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::init_value

The init_value element for ScanOpT.

Definition at line 191 of file agent_scan.cuh.

◆ scan_op

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >

ScanOpT cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::scan_op

Binary scan operator.

Definition at line 190 of file agent_scan.cuh.

◆ temp_storage

template<typename AgentScanPolicyT , typename InputIteratorT , typename OutputIteratorT , typename ScanOpT , typename InitValueT , typename OffsetT >

_TempStorage& cub::AgentScan< AgentScanPolicyT, InputIteratorT, OutputIteratorT, ScanOpT, InitValueT, OffsetT >::temp_storage

Reference to temp_storage.

Definition at line 187 of file agent_scan.cuh.

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

openfpm_data/src/util/cuda/cub_old/agent/agent_scan.cuh

Detailed Description

Data Structures

Public Types

Public Member Functions

Data Fields

Member Typedef Documentation

◆ BlockLoadT

◆ BlockScanT

◆ BlockStoreT

◆ InputT

◆ OutputT

◆ RunningPrefixCallbackOp

◆ ScanTileStateT

◆ TilePrefixCallbackOpT

◆ WrappedInputIteratorT

Member Enumeration Documentation

◆ anonymous enum

Constructor & Destructor Documentation

◆ AgentScan()

Member Function Documentation

◆ ConsumeRange() [1/3]

◆ ConsumeRange() [2/3]

◆ ConsumeRange() [3/3]

◆ ConsumeTile() [1/2]

◆ ConsumeTile() [2/2]

◆ ScanTile() [1/4]

◆ ScanTile() [2/4]

◆ ScanTile() [3/4]

◆ ScanTile() [4/4]

Field Documentation

◆ d_in

◆ d_out

◆ init_value

◆ scan_op

◆ temp_storage