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

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.

Member Typedef Documentation

◆ BlockDiscontinuityT

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

typedef BlockDiscontinuity< OutputT, BLOCK_THREADS> cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::BlockDiscontinuityT

Definition at line 176 of file agent_select_if.cuh.

◆ BlockLoadFlags

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

typedef BlockLoad< FlagT, BLOCK_THREADS, ITEMS_PER_THREAD, AgentSelectIfPolicyT::LOAD_ALGORITHM> cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::BlockLoadFlags

Definition at line 170 of file agent_select_if.cuh.

◆ BlockLoadT

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

typedef BlockLoad< OutputT, BLOCK_THREADS, ITEMS_PER_THREAD, AgentSelectIfPolicyT::LOAD_ALGORITHM> cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::BlockLoadT

Definition at line 162 of file agent_select_if.cuh.

◆ BlockScanT

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

typedef BlockScan< OffsetT, BLOCK_THREADS, AgentSelectIfPolicyT::SCAN_ALGORITHM> cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::BlockScanT

Definition at line 183 of file agent_select_if.cuh.

◆ FlagT

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

typedef std::iterator_traits<FlagsInputIteratorT>::value_type cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::FlagT

Definition at line 120 of file agent_select_if.cuh.

◆ InputT

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

typedef std::iterator_traits<InputIteratorT>::value_type cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::InputT

Definition at line 112 of file agent_select_if.cuh.

◆ ItemExchangeT

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

typedef OutputT cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::ItemExchangeT[TILE_ITEMS]

Definition at line 193 of file agent_select_if.cuh.

◆ OutputT

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

typedef If<(Equals<typenamestd::iterator_traits<SelectedOutputIteratorT>::value_type,void>::VALUE),typenamestd::iterator_traits<InputIteratorT>::value_type,typenamestd::iterator_traits<SelectedOutputIteratorT>::value_type>::Type cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::OutputT

Definition at line 117 of file agent_select_if.cuh.

◆ ScanTileStateT

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

typedef ScanTileState<OffsetT> cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::ScanTileStateT

Definition at line 123 of file agent_select_if.cuh.

◆ TilePrefixCallbackOpT

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

typedef TilePrefixCallbackOp< OffsetT, cub::Sum, ScanTileStateT> cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::TilePrefixCallbackOpT

Definition at line 190 of file agent_select_if.cuh.

◆ WrappedFlagsInputIteratorT

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

typedef If<IsPointer<FlagsInputIteratorT>::VALUE,CacheModifiedInputIterator<AgentSelectIfPolicyT::LOAD_MODIFIER,FlagT,OffsetT>,FlagsInputIteratorT>::Type cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::WrappedFlagsInputIteratorT

Definition at line 154 of file agent_select_if.cuh.

◆ WrappedInputIteratorT

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

typedef If<IsPointer<InputIteratorT>::VALUE,CacheModifiedInputIterator<AgentSelectIfPolicyT::LOAD_MODIFIER,InputT,OffsetT>,InputIteratorT>::Type cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::WrappedInputIteratorT

Definition at line 148 of file agent_select_if.cuh.

Member Enumeration Documentation

◆ anonymous enum

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

anonymous enum

Definition at line 126 of file agent_select_if.cuh.

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

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

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

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

Field Documentation

◆ d_flags_in

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

WrappedFlagsInputIteratorT cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::d_flags_in

Input selection flags (if applicable)

Definition at line 226 of file agent_select_if.cuh.

◆ d_in

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

WrappedInputIteratorT cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::d_in

Input items.

Definition at line 224 of file agent_select_if.cuh.

◆ d_selected_out

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

SelectedOutputIteratorT cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::d_selected_out

Unique output items.

Definition at line 225 of file agent_select_if.cuh.

◆ inequality_op

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

InequalityWrapper<EqualityOpT> cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::inequality_op

T inequality operator.

Definition at line 227 of file agent_select_if.cuh.

◆ num_items

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

OffsetT cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::num_items

Total number of input items.

Definition at line 229 of file agent_select_if.cuh.

◆ select_op

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

SelectOpT cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::select_op

Selection operator.

Definition at line 228 of file agent_select_if.cuh.

◆ temp_storage

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

_TempStorage& cub::AgentSelectIf< AgentSelectIfPolicyT, InputIteratorT, FlagsInputIteratorT, SelectedOutputIteratorT, SelectOpT, EqualityOpT, OffsetT, KEEP_REJECTS >::temp_storage

Reference to temp_storage.

Definition at line 223 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

Member Typedef Documentation

◆ BlockDiscontinuityT

◆ BlockLoadFlags

◆ BlockLoadT

◆ BlockScanT

◆ FlagT

◆ InputT

◆ ItemExchangeT

◆ OutputT

◆ ScanTileStateT

◆ TilePrefixCallbackOpT

◆ WrappedFlagsInputIteratorT

◆ WrappedInputIteratorT

Member Enumeration Documentation

◆ anonymous enum

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]

Field Documentation

◆ d_flags_in

◆ d_in

◆ d_selected_out

◆ inequality_op

◆ num_items

◆ select_op

◆ temp_storage