AgentSelectIf implements a stateful abstraction of CUDA thread blocks for participating in device-wide selection. More...

Detailed Description

template<typename AgentSelectIfPolicyT, typename InputIteratorT, typename FlagsInputIteratorT, typename SelectedOutputIteratorT, typename SelectOpT, typename EqualityOpT, typename OffsetT, bool KEEP_REJECTS>
struct cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >

AgentSelectIf implements a stateful abstraction of CUDA thread blocks for participating in device-wide selection.

Performs functor-based selection if SelectOpT functor type != NullType Otherwise performs flag-based selection if FlagsInputIterator's value type != NullType Otherwise performs discontinuity selection (keep unique)< Whether or not we push rejected items to the back of the output

Definition at line 105 of file agent_select_if.cuh.

Data Structures
union	_TempStorage

struct	TempStorage

Public Types
enum	{ USE_SELECT_OP, USE_SELECT_FLAGS, USE_DISCONTINUITY, BLOCK_THREADS = AgentSelectIfPolicyT::BLOCK_THREADS, ITEMS_PER_THREAD = AgentSelectIfPolicyT::ITEMS_PER_THREAD, TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD, TWO_PHASE_SCATTER = (ITEMS_PER_THREAD > 1), SELECT_METHOD }

typedef std::iterator_traits< InputIteratorT >::value_type	InputT

typedef If<(Equals< typename std::iterator_traits< SelectedOutputIteratorT >::value_type, void >::VALUE), typename std::iterator_traits< InputIteratorT >::value_type, typename std::iterator_traits< SelectedOutputIteratorT >::value_type >::Type	OutputT

typedef std::iterator_traits< FlagsInputIteratorT >::value_type	FlagT

typedef ScanTileState< OffsetT >	ScanTileStateT

typedef If< IsPointer< InputIteratorT >::VALUE, CacheModifiedInputIterator< AgentSelectIfPolicyT::LOAD_MODIFIER, InputT, OffsetT >, InputIteratorT >::Type	WrappedInputIteratorT

typedef If< IsPointer< FlagsInputIteratorT >::VALUE, CacheModifiedInputIterator< AgentSelectIfPolicyT::LOAD_MODIFIER, FlagT, OffsetT >, FlagsInputIteratorT >::Type	WrappedFlagsInputIteratorT

typedef BlockLoad< OutputT, BLOCK_THREADS, ITEMS_PER_THREAD, AgentSelectIfPolicyT::LOAD_ALGORITHM >	BlockLoadT

typedef BlockLoad< FlagT, BLOCK_THREADS, ITEMS_PER_THREAD, AgentSelectIfPolicyT::LOAD_ALGORITHM >	BlockLoadFlags

typedef BlockDiscontinuity< OutputT, BLOCK_THREADS >	BlockDiscontinuityT

typedef BlockScan< OffsetT, BLOCK_THREADS, AgentSelectIfPolicyT::SCAN_ALGORITHM >	BlockScanT

typedef TilePrefixCallbackOp< OffsetT, cub::Sum, ScanTileStateT >	TilePrefixCallbackOpT

typedef OutputT	ItemExchangeT[TILE_ITEMS]

Public Member Functions
__device__ __forceinline__	AgentSelectIf (TempStorage &temp_storage, InputIteratorT d_in, FlagsInputIteratorT d_flags_in, SelectedOutputIteratorT d_selected_out, SelectOpT select_op, EqualityOpT equality_op, OffsetT num_items)

template<bool IS_FIRST_TILE, bool IS_LAST_TILE>
__device__ __forceinline__ void	InitializeSelections (OffsetT, OffsetT num_tile_items, OutputT(&items)[ITEMS_PER_THREAD], OffsetT(&selection_flags)[ITEMS_PER_THREAD], Int2Type< USE_SELECT_OP >)

template<bool IS_FIRST_TILE, bool IS_LAST_TILE>
__device__ __forceinline__ void	InitializeSelections (OffsetT tile_offset, OffsetT num_tile_items, OutputT(&)[ITEMS_PER_THREAD], OffsetT(&selection_flags)[ITEMS_PER_THREAD], Int2Type< USE_SELECT_FLAGS >)

template<bool IS_FIRST_TILE, bool IS_LAST_TILE>
__device__ __forceinline__ void	InitializeSelections (OffsetT tile_offset, OffsetT num_tile_items, OutputT(&items)[ITEMS_PER_THREAD], OffsetT(&selection_flags)[ITEMS_PER_THREAD], Int2Type< USE_DISCONTINUITY >)

template<bool IS_LAST_TILE, bool IS_FIRST_TILE>
__device__ __forceinline__ void	ScatterDirect (OutputT(&items)[ITEMS_PER_THREAD], OffsetT(&selection_flags)[ITEMS_PER_THREAD], OffsetT(&selection_indices)[ITEMS_PER_THREAD], OffsetT num_selections)

template<bool IS_LAST_TILE, bool IS_FIRST_TILE>
__device__ __forceinline__ void	ScatterTwoPhase (OutputT(&items)[ITEMS_PER_THREAD], OffsetT(&selection_flags)[ITEMS_PER_THREAD], OffsetT(&selection_indices)[ITEMS_PER_THREAD], int, int num_tile_selections, OffsetT num_selections_prefix, OffsetT, Int2Type< false >)

template<bool IS_LAST_TILE, bool IS_FIRST_TILE>
__device__ __forceinline__ void	ScatterTwoPhase (OutputT(&items)[ITEMS_PER_THREAD], OffsetT(&selection_flags)[ITEMS_PER_THREAD], OffsetT(&selection_indices)[ITEMS_PER_THREAD], int num_tile_items, int num_tile_selections, OffsetT num_selections_prefix, OffsetT num_rejected_prefix, Int2Type< true >)

template<bool IS_LAST_TILE, bool IS_FIRST_TILE>
__device__ __forceinline__ void	Scatter (OutputT(&items)[ITEMS_PER_THREAD], OffsetT(&selection_flags)[ITEMS_PER_THREAD], OffsetT(&selection_indices)[ITEMS_PER_THREAD], int num_tile_items, int num_tile_selections, OffsetT num_selections_prefix, OffsetT num_rejected_prefix, OffsetT num_selections)

template<bool IS_LAST_TILE>
__device__ __forceinline__ OffsetT	ConsumeFirstTile (int num_tile_items, OffsetT tile_offset, ScanTileStateT &tile_state)

template<bool IS_LAST_TILE>
__device__ __forceinline__ OffsetT	ConsumeSubsequentTile (int num_tile_items, int tile_idx, OffsetT tile_offset, ScanTileStateT &tile_state)

template<bool IS_LAST_TILE>
__device__ __forceinline__ OffsetT	ConsumeTile (int num_tile_items, int tile_idx, OffsetT tile_offset, ScanTileStateT &tile_state)

template<typename NumSelectedIteratorT >
__device__ __forceinline__ void	ConsumeRange (int num_tiles, ScanTileStateT &tile_state, NumSelectedIteratorT d_num_selected_out)
	< Output iterator type for recording number of items selection_flags More...

Data Fields
_TempStorage &	temp_storage
	Reference to temp_storage.

WrappedInputIteratorT	d_in
	Input items.

SelectedOutputIteratorT	d_selected_out
	Unique output items.

WrappedFlagsInputIteratorT	d_flags_in
	Input selection flags (if applicable)

InequalityWrapper< EqualityOpT >	inequality_op
	T inequality operator.

SelectOpT	select_op
	Selection operator.

OffsetT	num_items
	Total number of input items.

Constructor & Destructor Documentation

◆ AgentSelectIf()

template<typename AgentSelectIfPolicyT , typename InputIteratorT , typename FlagsInputIteratorT , typename SelectedOutputIteratorT , typename SelectOpT , typename EqualityOpT , typename OffsetT , bool KEEP_REJECTS>

__device__ __forceinline__ cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::AgentSelectIf	(	TempStorage &	temp_storage,
		InputIteratorT	d_in,
		FlagsInputIteratorT	d_flags_in,
		SelectedOutputIteratorT	d_selected_out,
		SelectOpT	select_op,
		EqualityOpT	equality_op,
		OffsetT	num_items
	)

inline

Parameters

temp_storage	Reference to temp_storage
d_in	Input data
d_flags_in	Input selection flags (if applicable)
d_selected_out	Output data
select_op	Selection operator
equality_op	Equality operator
num_items	Total number of input items

Definition at line 238 of file agent_select_if.cuh.

Member Function Documentation

◆ ConsumeFirstTile()

template<typename AgentSelectIfPolicyT , typename InputIteratorT , typename FlagsInputIteratorT , typename SelectedOutputIteratorT , typename SelectOpT , typename EqualityOpT , typename OffsetT , bool KEEP_REJECTS>

template<bool IS_LAST_TILE>

__device__ __forceinline__ OffsetT cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::ConsumeFirstTile	(	int	num_tile_items,
		OffsetT	tile_offset,
		ScanTileStateT &	tile_state
	)

inline

Process first tile of input (dynamic chained scan). Returns the running count of selections (including this tile)

Parameters

num_tile_items	Number of input items comprising this tile
tile_offset	Tile offset
tile_state	Global tile state descriptor

Definition at line 524 of file agent_select_if.cuh.

◆ ConsumeRange()

template<typename AgentSelectIfPolicyT , typename InputIteratorT , typename FlagsInputIteratorT , typename SelectedOutputIteratorT , typename SelectOpT , typename EqualityOpT , typename OffsetT , bool KEEP_REJECTS>

template<typename NumSelectedIteratorT >

__device__ __forceinline__ void cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::ConsumeRange	(	int	num_tiles,
		ScanTileStateT &	tile_state,
		NumSelectedIteratorT	d_num_selected_out
	)

inline

< Output iterator type for recording number of items selection_flags

Scan tiles of items as part of a dynamic chained scan

Parameters

num_tiles	Total number of input tiles
tile_state	Global tile state descriptor
d_num_selected_out	Output total number selection_flags

Definition at line 669 of file agent_select_if.cuh.

◆ ConsumeSubsequentTile()

template<typename AgentSelectIfPolicyT , typename InputIteratorT , typename FlagsInputIteratorT , typename SelectedOutputIteratorT , typename SelectOpT , typename EqualityOpT , typename OffsetT , bool KEEP_REJECTS>

template<bool IS_LAST_TILE>

__device__ __forceinline__ OffsetT cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::ConsumeSubsequentTile	(	int	num_tile_items,
		int	tile_idx,
		OffsetT	tile_offset,
		ScanTileStateT &	tile_state
	)

inline

Process subsequent tile of input (dynamic chained scan). Returns the running count of selections (including this tile)

Parameters

num_tile_items	Number of input items comprising this tile
tile_idx	Tile index
tile_offset	Tile offset
tile_state	Global tile state descriptor

Definition at line 583 of file agent_select_if.cuh.

◆ ConsumeTile()

template<typename AgentSelectIfPolicyT , typename InputIteratorT , typename FlagsInputIteratorT , typename SelectedOutputIteratorT , typename SelectOpT , typename EqualityOpT , typename OffsetT , bool KEEP_REJECTS>

template<bool IS_LAST_TILE>

__device__ __forceinline__ OffsetT cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::ConsumeTile	(	int	num_tile_items,
		int	tile_idx,
		OffsetT	tile_offset,
		ScanTileStateT &	tile_state
	)

inline

Process a tile of input

Parameters

num_tile_items	Number of input items comprising this tile
tile_idx	Tile index
tile_offset	Tile offset
tile_state	Global tile state descriptor

Definition at line 645 of file agent_select_if.cuh.

◆ InitializeSelections() [1/3]

template<typename AgentSelectIfPolicyT , typename InputIteratorT , typename FlagsInputIteratorT , typename SelectedOutputIteratorT , typename SelectOpT , typename EqualityOpT , typename OffsetT , bool KEEP_REJECTS>

template<bool IS_FIRST_TILE, bool IS_LAST_TILE>

__device__ __forceinline__ void cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::InitializeSelections	(	OffsetT	,
		OffsetT	num_tile_items,
		OutputT(&)	items[ITEMS_PER_THREAD],
		OffsetT(&)	selection_flags[ITEMS_PER_THREAD],
		Int2Type< USE_SELECT_OP >
	)

inline

Initialize selections (specialized for selection operator)

Definition at line 265 of file agent_select_if.cuh.

◆ InitializeSelections() [2/3]

template<typename AgentSelectIfPolicyT , typename InputIteratorT , typename FlagsInputIteratorT , typename SelectedOutputIteratorT , typename SelectOpT , typename EqualityOpT , typename OffsetT , bool KEEP_REJECTS>

template<bool IS_FIRST_TILE, bool IS_LAST_TILE>

__device__ __forceinline__ void cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::InitializeSelections	(	OffsetT	tile_offset,
		OffsetT	num_tile_items,
		OutputT(&)	[ITEMS_PER_THREAD],
		OffsetT(&)	selection_flags[ITEMS_PER_THREAD],
		Int2Type< USE_SELECT_FLAGS >
	)

inline

Initialize selections (specialized for valid flags)

Definition at line 288 of file agent_select_if.cuh.

◆ InitializeSelections() [3/3]

template<typename AgentSelectIfPolicyT , typename InputIteratorT , typename FlagsInputIteratorT , typename SelectedOutputIteratorT , typename SelectOpT , typename EqualityOpT , typename OffsetT , bool KEEP_REJECTS>

template<bool IS_FIRST_TILE, bool IS_LAST_TILE>

__device__ __forceinline__ void cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::InitializeSelections	(	OffsetT	tile_offset,
		OffsetT	num_tile_items,
		OutputT(&)	items[ITEMS_PER_THREAD],
		OffsetT(&)	selection_flags[ITEMS_PER_THREAD],
		Int2Type< USE_DISCONTINUITY >
	)

inline

Initialize selections (specialized for discontinuity detection)

Definition at line 322 of file agent_select_if.cuh.

◆ Scatter()

template<typename AgentSelectIfPolicyT , typename InputIteratorT , typename FlagsInputIteratorT , typename SelectedOutputIteratorT , typename SelectOpT , typename EqualityOpT , typename OffsetT , bool KEEP_REJECTS>

template<bool IS_LAST_TILE, bool IS_FIRST_TILE>

__device__ __forceinline__ void cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::Scatter	(	OutputT(&)	items[ITEMS_PER_THREAD],
		OffsetT(&)	selection_flags[ITEMS_PER_THREAD],
		OffsetT(&)	selection_indices[ITEMS_PER_THREAD],
		int	num_tile_items,
		int	num_tile_selections,
		OffsetT	num_selections_prefix,
		OffsetT	num_rejected_prefix,
		OffsetT	num_selections
	)

inline

Scatter flagged items

Parameters

num_tile_items	Number of valid items in this tile
num_tile_selections	Number of selections in this tile
num_selections_prefix	Total number of selections prior to this tile
num_rejected_prefix	Total number of rejections prior to this tile
num_selections	Total number of selections including this tile

Definition at line 482 of file agent_select_if.cuh.

◆ ScatterDirect()

template<typename AgentSelectIfPolicyT , typename InputIteratorT , typename FlagsInputIteratorT , typename SelectedOutputIteratorT , typename SelectOpT , typename EqualityOpT , typename OffsetT , bool KEEP_REJECTS>

template<bool IS_LAST_TILE, bool IS_FIRST_TILE>

__device__ __forceinline__ void cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::ScatterDirect	(	OutputT(&)	items[ITEMS_PER_THREAD],
		OffsetT(&)	selection_flags[ITEMS_PER_THREAD],
		OffsetT(&)	selection_indices[ITEMS_PER_THREAD],
		OffsetT	num_selections
	)

inline

Scatter flagged items to output offsets (specialized for direct scattering)

Definition at line 366 of file agent_select_if.cuh.

◆ ScatterTwoPhase() [1/2]

template<typename AgentSelectIfPolicyT , typename InputIteratorT , typename FlagsInputIteratorT , typename SelectedOutputIteratorT , typename SelectOpT , typename EqualityOpT , typename OffsetT , bool KEEP_REJECTS>

template<bool IS_LAST_TILE, bool IS_FIRST_TILE>

__device__ __forceinline__ void cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::ScatterTwoPhase	(	OutputT(&)	items[ITEMS_PER_THREAD],
		OffsetT(&)	selection_flags[ITEMS_PER_THREAD],
		OffsetT(&)	selection_indices[ITEMS_PER_THREAD],
		int	,
		int	num_tile_selections,
		OffsetT	num_selections_prefix,
		OffsetT	,
		Int2Type< false >
	)

inline

Scatter flagged items to output offsets (specialized for two-phase scattering)

Parameters

num_tile_selections	Number of selections in this tile
num_selections_prefix	Total number of selections prior to this tile

Definition at line 391 of file agent_select_if.cuh.

◆ ScatterTwoPhase() [2/2]

template<typename AgentSelectIfPolicyT , typename InputIteratorT , typename FlagsInputIteratorT , typename SelectedOutputIteratorT , typename SelectOpT , typename EqualityOpT , typename OffsetT , bool KEEP_REJECTS>

template<bool IS_LAST_TILE, bool IS_FIRST_TILE>

__device__ __forceinline__ void cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::ScatterTwoPhase	(	OutputT(&)	items[ITEMS_PER_THREAD],
		OffsetT(&)	selection_flags[ITEMS_PER_THREAD],
		OffsetT(&)	selection_indices[ITEMS_PER_THREAD],
		int	num_tile_items,
		int	num_tile_selections,
		OffsetT	num_selections_prefix,
		OffsetT	num_rejected_prefix,
		Int2Type< true >
	)

inline

Scatter flagged items to output offsets (specialized for two-phase scattering)

Parameters

num_tile_items	Number of valid items in this tile
num_tile_selections	Number of selections in this tile
num_selections_prefix	Total number of selections prior to this tile
num_rejected_prefix	Total number of rejections prior to this tile

Definition at line 427 of file agent_select_if.cuh.

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

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

Detailed Description

Data Structures

Public Types

Public Member Functions

Data Fields

Constructor & Destructor Documentation

◆ AgentSelectIf()

Member Function Documentation

◆ ConsumeFirstTile()

◆ ConsumeRange()

◆ ConsumeSubsequentTile()

◆ ConsumeTile()

◆ InitializeSelections() [1/3]

◆ InitializeSelections() [2/3]

◆ InitializeSelections() [3/3]

◆ Scatter()

◆ ScatterDirect()

◆ ScatterTwoPhase() [1/2]

◆ ScatterTwoPhase() [2/2]