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

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],
		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:

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

Detailed Description

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

Data Structures

Public Types

Public Member Functions

Data Fields

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]

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