AgentReduceByKey implements a stateful abstraction of CUDA thread blocks for participating in device-wide reduce-value-by-key. More...

Detailed Description

template<typename AgentReduceByKeyPolicyT, typename KeysInputIteratorT, typename UniqueOutputIteratorT, typename ValuesInputIteratorT, typename AggregatesOutputIteratorT, typename NumRunsOutputIteratorT, typename EqualityOpT, typename ReductionOpT, typename OffsetT>
struct cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >

AgentReduceByKey implements a stateful abstraction of CUDA thread blocks for participating in device-wide reduce-value-by-key.

< Signed integer type for global offsets

Definition at line 98 of file agent_reduce_by_key.cuh.

Data Structures
union	_TempStorage

struct	GuardedInequalityWrapper

struct	TempStorage

Public Types
enum	{ BLOCK_THREADS = AgentReduceByKeyPolicyT::BLOCK_THREADS , ITEMS_PER_THREAD = AgentReduceByKeyPolicyT::ITEMS_PER_THREAD , TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD , TWO_PHASE_SCATTER = (ITEMS_PER_THREAD > 1) , HAS_IDENTITY_ZERO = (Equals<ReductionOpT, cub::Sum>::VALUE) && (Traits<ValueOutputT>::PRIMITIVE) }

typedef std::iterator_traits< KeysInputIteratorT >::value_type	KeyInputT

typedef If<(Equals< typenamestd::iterator_traits< UniqueOutputIteratorT >::value_type, void >::VALUE), typenamestd::iterator_traits< KeysInputIteratorT >::value_type, typenamestd::iterator_traits< UniqueOutputIteratorT >::value_type >::Type	KeyOutputT

typedef std::iterator_traits< ValuesInputIteratorT >::value_type	ValueInputT

typedef If<(Equals< typenamestd::iterator_traits< AggregatesOutputIteratorT >::value_type, void >::VALUE), typenamestd::iterator_traits< ValuesInputIteratorT >::value_type, typenamestd::iterator_traits< AggregatesOutputIteratorT >::value_type >::Type	ValueOutputT

typedef KeyValuePair< OffsetT, ValueOutputT >	OffsetValuePairT

typedef KeyValuePair< KeyOutputT, ValueOutputT >	KeyValuePairT

typedef ReduceByKeyScanTileState< ValueOutputT, OffsetT >	ScanTileStateT

typedef If< IsPointer< KeysInputIteratorT >::VALUE, CacheModifiedInputIterator< AgentReduceByKeyPolicyT::LOAD_MODIFIER, KeyInputT, OffsetT >, KeysInputIteratorT >::Type	WrappedKeysInputIteratorT

typedef If< IsPointer< ValuesInputIteratorT >::VALUE, CacheModifiedInputIterator< AgentReduceByKeyPolicyT::LOAD_MODIFIER, ValueInputT, OffsetT >, ValuesInputIteratorT >::Type	WrappedValuesInputIteratorT

typedef If< IsPointer< AggregatesOutputIteratorT >::VALUE, CacheModifiedInputIterator< AgentReduceByKeyPolicyT::LOAD_MODIFIER, ValueInputT, OffsetT >, AggregatesOutputIteratorT >::Type	WrappedFixupInputIteratorT

typedef ReduceBySegmentOp< ReductionOpT >	ReduceBySegmentOpT

typedef BlockLoad< KeyOutputT, BLOCK_THREADS, ITEMS_PER_THREAD, AgentReduceByKeyPolicyT::LOAD_ALGORITHM >	BlockLoadKeysT

typedef BlockLoad< ValueOutputT, BLOCK_THREADS, ITEMS_PER_THREAD, AgentReduceByKeyPolicyT::LOAD_ALGORITHM >	BlockLoadValuesT

typedef BlockDiscontinuity< KeyOutputT, BLOCK_THREADS >	BlockDiscontinuityKeys

typedef BlockScan< OffsetValuePairT, BLOCK_THREADS, AgentReduceByKeyPolicyT::SCAN_ALGORITHM >	BlockScanT

typedef TilePrefixCallbackOp< OffsetValuePairT, ReduceBySegmentOpT, ScanTileStateT >	TilePrefixCallbackOpT

typedef KeyOutputT	KeyExchangeT[TILE_ITEMS+1]

typedef ValueOutputT	ValueExchangeT[TILE_ITEMS+1]

Public Member Functions
__device__ __forceinline__	AgentReduceByKey (TempStorage &temp_storage, KeysInputIteratorT d_keys_in, UniqueOutputIteratorT d_unique_out, ValuesInputIteratorT d_values_in, AggregatesOutputIteratorT d_aggregates_out, NumRunsOutputIteratorT d_num_runs_out, EqualityOpT equality_op, ReductionOpT reduction_op)

__device__ __forceinline__ void	ScatterDirect (KeyValuePairT(&scatter_items)[ITEMS_PER_THREAD], OffsetT(&segment_flags)[ITEMS_PER_THREAD], OffsetT(&segment_indices)[ITEMS_PER_THREAD])

__device__ __forceinline__ void	ScatterTwoPhase (KeyValuePairT(&scatter_items)[ITEMS_PER_THREAD], OffsetT(&segment_flags)[ITEMS_PER_THREAD], OffsetT(&segment_indices)[ITEMS_PER_THREAD], OffsetT num_tile_segments, OffsetT num_tile_segments_prefix)

__device__ __forceinline__ void	Scatter (KeyValuePairT(&scatter_items)[ITEMS_PER_THREAD], OffsetT(&segment_flags)[ITEMS_PER_THREAD], OffsetT(&segment_indices)[ITEMS_PER_THREAD], OffsetT num_tile_segments, OffsetT num_tile_segments_prefix)

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)

Data Fields
_TempStorage &	temp_storage
	Reference to temp_storage.

WrappedKeysInputIteratorT	d_keys_in
	Input keys.

UniqueOutputIteratorT	d_unique_out
	Unique output keys.

WrappedValuesInputIteratorT	d_values_in
	Input values.

AggregatesOutputIteratorT	d_aggregates_out
	Output value aggregates.

NumRunsOutputIteratorT	d_num_runs_out
	Output pointer for total number of segments identified.

EqualityOpT	equality_op
	KeyT equality operator.

ReductionOpT	reduction_op
	Reduction operator.

ReduceBySegmentOpT	scan_op
	Reduce-by-segment scan operator.

Member Typedef Documentation

◆ BlockDiscontinuityKeys

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef BlockDiscontinuity< KeyOutputT, BLOCK_THREADS> cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::BlockDiscontinuityKeys

Definition at line 206 of file agent_reduce_by_key.cuh.

◆ BlockLoadKeysT

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef BlockLoad< KeyOutputT, BLOCK_THREADS, ITEMS_PER_THREAD, AgentReduceByKeyPolicyT::LOAD_ALGORITHM> cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::BlockLoadKeysT

Definition at line 192 of file agent_reduce_by_key.cuh.

◆ BlockLoadValuesT

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef BlockLoad< ValueOutputT, BLOCK_THREADS, ITEMS_PER_THREAD, AgentReduceByKeyPolicyT::LOAD_ALGORITHM> cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::BlockLoadValuesT

Definition at line 200 of file agent_reduce_by_key.cuh.

◆ BlockScanT

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef BlockScan< OffsetValuePairT, BLOCK_THREADS, AgentReduceByKeyPolicyT::SCAN_ALGORITHM> cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::BlockScanT

Definition at line 213 of file agent_reduce_by_key.cuh.

◆ KeyExchangeT

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef KeyOutputT cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::KeyExchangeT[TILE_ITEMS+1]

Definition at line 223 of file agent_reduce_by_key.cuh.

◆ KeyInputT

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef std::iterator_traits<KeysInputIteratorT>::value_type cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::KeyInputT

Definition at line 105 of file agent_reduce_by_key.cuh.

◆ KeyOutputT

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef If<(Equals<typenamestd::iterator_traits<UniqueOutputIteratorT>::value_type,void>::VALUE),typenamestd::iterator_traits<KeysInputIteratorT>::value_type,typenamestd::iterator_traits<UniqueOutputIteratorT>::value_type>::Type cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::KeyOutputT

Definition at line 110 of file agent_reduce_by_key.cuh.

◆ KeyValuePairT

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef KeyValuePair<KeyOutputT, ValueOutputT> cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::KeyValuePairT

Definition at line 124 of file agent_reduce_by_key.cuh.

◆ OffsetValuePairT

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef KeyValuePair<OffsetT, ValueOutputT> cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::OffsetValuePairT

Definition at line 121 of file agent_reduce_by_key.cuh.

◆ ReduceBySegmentOpT

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef ReduceBySegmentOp<ReductionOpT> cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::ReduceBySegmentOpT

Definition at line 184 of file agent_reduce_by_key.cuh.

◆ ScanTileStateT

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef ReduceByKeyScanTileState<ValueOutputT, OffsetT> cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::ScanTileStateT

Definition at line 127 of file agent_reduce_by_key.cuh.

◆ TilePrefixCallbackOpT

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef TilePrefixCallbackOp< OffsetValuePairT, ReduceBySegmentOpT, ScanTileStateT> cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::TilePrefixCallbackOpT

Definition at line 220 of file agent_reduce_by_key.cuh.

◆ ValueExchangeT

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef ValueOutputT cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::ValueExchangeT[TILE_ITEMS+1]

Definition at line 224 of file agent_reduce_by_key.cuh.

◆ ValueInputT

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef std::iterator_traits<ValuesInputIteratorT>::value_type cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::ValueInputT

Definition at line 113 of file agent_reduce_by_key.cuh.

◆ ValueOutputT

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef If<(Equals<typenamestd::iterator_traits<AggregatesOutputIteratorT>::value_type,void>::VALUE),typenamestd::iterator_traits<ValuesInputIteratorT>::value_type,typenamestd::iterator_traits<AggregatesOutputIteratorT>::value_type>::Type cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::ValueOutputT

Definition at line 118 of file agent_reduce_by_key.cuh.

◆ WrappedFixupInputIteratorT

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef If<IsPointer<AggregatesOutputIteratorT>::VALUE,CacheModifiedInputIterator<AgentReduceByKeyPolicyT::LOAD_MODIFIER,ValueInputT,OffsetT>,AggregatesOutputIteratorT>::Type cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::WrappedFixupInputIteratorT

Definition at line 181 of file agent_reduce_by_key.cuh.

◆ WrappedKeysInputIteratorT

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef If<IsPointer<KeysInputIteratorT>::VALUE,CacheModifiedInputIterator<AgentReduceByKeyPolicyT::LOAD_MODIFIER,KeyInputT,OffsetT>,KeysInputIteratorT>::Type cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::WrappedKeysInputIteratorT

Definition at line 169 of file agent_reduce_by_key.cuh.

◆ WrappedValuesInputIteratorT

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef If<IsPointer<ValuesInputIteratorT>::VALUE,CacheModifiedInputIterator<AgentReduceByKeyPolicyT::LOAD_MODIFIER,ValueInputT,OffsetT>,ValuesInputIteratorT>::Type cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::WrappedValuesInputIteratorT

Definition at line 175 of file agent_reduce_by_key.cuh.

Member Enumeration Documentation

◆ anonymous enum

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

anonymous enum

Definition at line 154 of file agent_reduce_by_key.cuh.

Constructor & Destructor Documentation

◆ AgentReduceByKey()

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

__device__ __forceinline__ cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::AgentReduceByKey	(	TempStorage &	temp_storage,
		KeysInputIteratorT	d_keys_in,
		UniqueOutputIteratorT	d_unique_out,
		ValuesInputIteratorT	d_values_in,
		AggregatesOutputIteratorT	d_aggregates_out,
		NumRunsOutputIteratorT	d_num_runs_out,
		EqualityOpT	equality_op,
		ReductionOpT	reduction_op
	)

inline

Parameters

temp_storage	Reference to temp_storage
d_keys_in	Input keys
d_unique_out	Unique output keys
d_values_in	Input values
d_aggregates_out	Output value aggregates
d_num_runs_out	Output pointer for total number of segments identified
equality_op	KeyT equality operator
reduction_op	ValueT reduction operator

Definition at line 271 of file agent_reduce_by_key.cuh.

Member Function Documentation

◆ ConsumeRange()

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

__device__ __forceinline__ void cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, 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 520 of file agent_reduce_by_key.cuh.

◆ ConsumeTile()

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

template<bool IS_LAST_TILE>

__device__ __forceinline__ void cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, 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 392 of file agent_reduce_by_key.cuh.

◆ Scatter()

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

__device__ __forceinline__ void cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::Scatter	(	KeyValuePairT(&)	scatter_items[ITEMS_PER_THREAD],
		OffsetT(&)	segment_flags[ITEMS_PER_THREAD],
		OffsetT(&)	segment_indices[ITEMS_PER_THREAD],
		OffsetT	num_tile_segments,
		OffsetT	num_tile_segments_prefix
	)

inline

Scatter flagged items

Definition at line 357 of file agent_reduce_by_key.cuh.

◆ ScatterDirect()

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

__device__ __forceinline__ void cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::ScatterDirect	(	KeyValuePairT(&)	scatter_items[ITEMS_PER_THREAD],
		OffsetT(&)	segment_flags[ITEMS_PER_THREAD],
		OffsetT(&)	segment_indices[ITEMS_PER_THREAD]
	)

inline

Directly scatter flagged items to output offsets

Definition at line 300 of file agent_reduce_by_key.cuh.

◆ ScatterTwoPhase()

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

__device__ __forceinline__ void cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::ScatterTwoPhase	(	KeyValuePairT(&)	scatter_items[ITEMS_PER_THREAD],
		OffsetT(&)	segment_flags[ITEMS_PER_THREAD],
		OffsetT(&)	segment_indices[ITEMS_PER_THREAD],
		OffsetT	num_tile_segments,
		OffsetT	num_tile_segments_prefix
	)

inline

2-phase scatter flagged items to output offsets

The exclusive scan causes each head flag to be paired with the previous value aggregate: the scatter offsets must be decremented for value aggregates

Definition at line 324 of file agent_reduce_by_key.cuh.

Field Documentation

◆ d_aggregates_out

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

AggregatesOutputIteratorT cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::d_aggregates_out

Output value aggregates.

Definition at line 258 of file agent_reduce_by_key.cuh.

◆ d_keys_in

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

WrappedKeysInputIteratorT cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::d_keys_in

Input keys.

Definition at line 255 of file agent_reduce_by_key.cuh.

◆ d_num_runs_out

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

NumRunsOutputIteratorT cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::d_num_runs_out

Output pointer for total number of segments identified.

Definition at line 259 of file agent_reduce_by_key.cuh.

◆ d_unique_out

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

UniqueOutputIteratorT cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::d_unique_out

Unique output keys.

Definition at line 256 of file agent_reduce_by_key.cuh.

◆ d_values_in

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

WrappedValuesInputIteratorT cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::d_values_in

Input values.

Definition at line 257 of file agent_reduce_by_key.cuh.

◆ equality_op

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

EqualityOpT cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::equality_op

KeyT equality operator.

Definition at line 260 of file agent_reduce_by_key.cuh.

◆ reduction_op

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

ReductionOpT cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::reduction_op

Reduction operator.

Definition at line 261 of file agent_reduce_by_key.cuh.

◆ scan_op

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

ReduceBySegmentOpT cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::scan_op

Reduce-by-segment scan operator.

Definition at line 262 of file agent_reduce_by_key.cuh.

◆ temp_storage

template<typename AgentReduceByKeyPolicyT , typename KeysInputIteratorT , typename UniqueOutputIteratorT , typename ValuesInputIteratorT , typename AggregatesOutputIteratorT , typename NumRunsOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

_TempStorage& cub::AgentReduceByKey< AgentReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, AggregatesOutputIteratorT, NumRunsOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::temp_storage

Reference to temp_storage.

Definition at line 254 of file agent_reduce_by_key.cuh.

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

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

Detailed Description

Data Structures

Public Types

Public Member Functions

Data Fields

Member Typedef Documentation

◆ BlockDiscontinuityKeys

◆ BlockLoadKeysT

◆ BlockLoadValuesT

◆ BlockScanT

◆ KeyExchangeT

◆ KeyInputT

◆ KeyOutputT

◆ KeyValuePairT

◆ OffsetValuePairT

◆ ReduceBySegmentOpT

◆ ScanTileStateT

◆ TilePrefixCallbackOpT

◆ ValueExchangeT

◆ ValueInputT

◆ ValueOutputT

◆ WrappedFixupInputIteratorT

◆ WrappedKeysInputIteratorT

◆ WrappedValuesInputIteratorT

Member Enumeration Documentation

◆ anonymous enum

Constructor & Destructor Documentation

◆ AgentReduceByKey()

Member Function Documentation

◆ ConsumeRange()

◆ ConsumeTile()

◆ Scatter()

◆ ScatterDirect()

◆ ScatterTwoPhase()

Field Documentation

◆ d_aggregates_out

◆ d_keys_in

◆ d_num_runs_out

◆ d_unique_out

◆ d_values_in

◆ equality_op

◆ reduction_op

◆ scan_op

◆ temp_storage