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

Detailed Description

template<typename AgentSegmentFixupPolicyT, typename PairsInputIteratorT, typename AggregatesOutputIteratorT, typename EqualityOpT, typename ReductionOpT, typename OffsetT>
struct cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >

AgentSegmentFixup 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 95 of file agent_segment_fixup.cuh.

Data Structures
union	_TempStorage

struct	TempStorage

Public Types
enum	{ BLOCK_THREADS = AgentSegmentFixupPolicyT::BLOCK_THREADS , ITEMS_PER_THREAD = AgentSegmentFixupPolicyT::ITEMS_PER_THREAD , TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD , USE_ATOMIC_FIXUP , HAS_IDENTITY_ZERO = (Equals<ReductionOpT, cub::Sum>::VALUE) && (Traits<ValueT>::PRIMITIVE) }

typedef std::iterator_traits< PairsInputIteratorT >::value_type	KeyValuePairT

typedef KeyValuePairT::Value	ValueT

typedef ReduceByKeyScanTileState< ValueT, OffsetT >	ScanTileStateT

typedef If< IsPointer< PairsInputIteratorT >::VALUE, CacheModifiedInputIterator< AgentSegmentFixupPolicyT::LOAD_MODIFIER, KeyValuePairT, OffsetT >, PairsInputIteratorT >::Type	WrappedPairsInputIteratorT

typedef If< IsPointer< AggregatesOutputIteratorT >::VALUE, CacheModifiedInputIterator< AgentSegmentFixupPolicyT::LOAD_MODIFIER, ValueT, OffsetT >, AggregatesOutputIteratorT >::Type	WrappedFixupInputIteratorT

typedef ReduceByKeyOp< cub::Sum >	ReduceBySegmentOpT

typedef BlockLoad< KeyValuePairT, BLOCK_THREADS, ITEMS_PER_THREAD, AgentSegmentFixupPolicyT::LOAD_ALGORITHM >	BlockLoadPairs

typedef BlockScan< KeyValuePairT, BLOCK_THREADS, AgentSegmentFixupPolicyT::SCAN_ALGORITHM >	BlockScanT

typedef TilePrefixCallbackOp< KeyValuePairT, ReduceBySegmentOpT, ScanTileStateT >	TilePrefixCallbackOpT

Public Member Functions
__device__ __forceinline__	AgentSegmentFixup (TempStorage &temp_storage, PairsInputIteratorT d_pairs_in, AggregatesOutputIteratorT d_aggregates_out, EqualityOpT equality_op, ReductionOpT reduction_op)

template<bool IS_LAST_TILE>
__device__ __forceinline__ void	ConsumeTile (OffsetT num_remaining, int tile_idx, OffsetT tile_offset, ScanTileStateT &tile_state, Int2Type< true > use_atomic_fixup)

template<bool IS_LAST_TILE>
__device__ __forceinline__ void	ConsumeTile (OffsetT num_remaining, int tile_idx, OffsetT tile_offset, ScanTileStateT &tile_state, Int2Type< false > use_atomic_fixup)

__device__ __forceinline__ void	ConsumeRange (int num_items, int num_tiles, ScanTileStateT &tile_state)

Data Fields
_TempStorage &	temp_storage
	Reference to temp_storage.

WrappedPairsInputIteratorT	d_pairs_in
	Input keys.

AggregatesOutputIteratorT	d_aggregates_out
	Output value aggregates.

WrappedFixupInputIteratorT	d_fixup_in
	Fixup input values.

InequalityWrapper< EqualityOpT >	inequality_op
	KeyT inequality operator.

ReductionOpT	reduction_op
	Reduction operator.

ReduceBySegmentOpT	scan_op
	Reduce-by-segment scan operator.

Member Typedef Documentation

◆ BlockLoadPairs

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef BlockLoad< KeyValuePairT, BLOCK_THREADS, ITEMS_PER_THREAD, AgentSegmentFixupPolicyT::LOAD_ALGORITHM> cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::BlockLoadPairs

Definition at line 149 of file agent_segment_fixup.cuh.

◆ BlockScanT

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef BlockScan< KeyValuePairT, BLOCK_THREADS, AgentSegmentFixupPolicyT::SCAN_ALGORITHM> cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::BlockScanT

Definition at line 156 of file agent_segment_fixup.cuh.

◆ KeyValuePairT

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef std::iterator_traits<PairsInputIteratorT>::value_type cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::KeyValuePairT

Definition at line 102 of file agent_segment_fixup.cuh.

◆ ReduceBySegmentOpT

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef ReduceByKeyOp<cub::Sum> cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::ReduceBySegmentOpT

Definition at line 141 of file agent_segment_fixup.cuh.

◆ ScanTileStateT

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef ReduceByKeyScanTileState<ValueT, OffsetT> cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::ScanTileStateT

Definition at line 108 of file agent_segment_fixup.cuh.

◆ TilePrefixCallbackOpT

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef TilePrefixCallbackOp< KeyValuePairT, ReduceBySegmentOpT, ScanTileStateT> cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::TilePrefixCallbackOpT

Definition at line 163 of file agent_segment_fixup.cuh.

◆ ValueT

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef KeyValuePairT::Value cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::ValueT

Definition at line 105 of file agent_segment_fixup.cuh.

◆ WrappedFixupInputIteratorT

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef If<IsPointer<AggregatesOutputIteratorT>::VALUE,CacheModifiedInputIterator<AgentSegmentFixupPolicyT::LOAD_MODIFIER,ValueT,OffsetT>,AggregatesOutputIteratorT>::Type cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::WrappedFixupInputIteratorT

Definition at line 138 of file agent_segment_fixup.cuh.

◆ WrappedPairsInputIteratorT

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

typedef If<IsPointer<PairsInputIteratorT>::VALUE,CacheModifiedInputIterator<AgentSegmentFixupPolicyT::LOAD_MODIFIER,KeyValuePairT,OffsetT>,PairsInputIteratorT>::Type cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::WrappedPairsInputIteratorT

Definition at line 132 of file agent_segment_fixup.cuh.

Member Enumeration Documentation

◆ anonymous enum

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

anonymous enum

Definition at line 111 of file agent_segment_fixup.cuh.

Constructor & Destructor Documentation

◆ AgentSegmentFixup()

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

__device__ __forceinline__ cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::AgentSegmentFixup	(	TempStorage &	temp_storage,
		PairsInputIteratorT	d_pairs_in,
		AggregatesOutputIteratorT	d_aggregates_out,
		EqualityOpT	equality_op,
		ReductionOpT	reduction_op
	)

inline

Parameters

temp_storage	Reference to temp_storage
d_pairs_in	Input keys
d_aggregates_out	Output value aggregates
equality_op	KeyT equality operator
reduction_op	ValueT reduction operator

Definition at line 201 of file agent_segment_fixup.cuh.

Member Function Documentation

◆ ConsumeRange()

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

__device__ __forceinline__ void cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::ConsumeRange	(	int	num_items,
		int	num_tiles,
		ScanTileStateT &	tile_state
	)

inline

Scan tiles of items as part of a dynamic chained scan

Parameters

num_items	Total number of input items
num_tiles	Total number of input tiles
tile_state	Global tile state descriptor

Definition at line 348 of file agent_segment_fixup.cuh.

◆ ConsumeTile() [1/2]

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

template<bool IS_LAST_TILE>

__device__ __forceinline__ void cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::ConsumeTile	(	OffsetT	num_remaining,
		int	tile_idx,
		OffsetT	tile_offset,
		ScanTileStateT &	tile_state,
		Int2Type< false >	use_atomic_fixup
	)

inline

Process input tile. Specialized for reduce-by-key fixup

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
use_atomic_fixup	Marker whether to use atomicAdd (instead of reduce-by-key)

Definition at line 267 of file agent_segment_fixup.cuh.

◆ ConsumeTile() [2/2]

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

template<bool IS_LAST_TILE>

__device__ __forceinline__ void cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::ConsumeTile	(	OffsetT	num_remaining,
		int	tile_idx,
		OffsetT	tile_offset,
		ScanTileStateT &	tile_state,
		Int2Type< true >	use_atomic_fixup
	)

inline

Process input tile. Specialized for atomic-fixup

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
use_atomic_fixup	Marker whether to use atomicAdd (instead of reduce-by-key)

Definition at line 227 of file agent_segment_fixup.cuh.

Field Documentation

◆ d_aggregates_out

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

AggregatesOutputIteratorT cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::d_aggregates_out

Output value aggregates.

Definition at line 188 of file agent_segment_fixup.cuh.

◆ d_fixup_in

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

WrappedFixupInputIteratorT cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::d_fixup_in

Fixup input values.

Definition at line 189 of file agent_segment_fixup.cuh.

◆ d_pairs_in

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

WrappedPairsInputIteratorT cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::d_pairs_in

Input keys.

Definition at line 187 of file agent_segment_fixup.cuh.

◆ inequality_op

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

InequalityWrapper<EqualityOpT> cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::inequality_op

KeyT inequality operator.

Definition at line 190 of file agent_segment_fixup.cuh.

◆ reduction_op

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

ReductionOpT cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::reduction_op

Reduction operator.

Definition at line 191 of file agent_segment_fixup.cuh.

◆ scan_op

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

ReduceBySegmentOpT cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::scan_op

Reduce-by-segment scan operator.

Definition at line 192 of file agent_segment_fixup.cuh.

◆ temp_storage

template<typename AgentSegmentFixupPolicyT , typename PairsInputIteratorT , typename AggregatesOutputIteratorT , typename EqualityOpT , typename ReductionOpT , typename OffsetT >

_TempStorage& cub::AgentSegmentFixup< AgentSegmentFixupPolicyT, PairsInputIteratorT, AggregatesOutputIteratorT, EqualityOpT, ReductionOpT, OffsetT >::temp_storage

Reference to temp_storage.

Definition at line 186 of file agent_segment_fixup.cuh.

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

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

Detailed Description

Data Structures

Public Types

Public Member Functions

Data Fields

Member Typedef Documentation

◆ BlockLoadPairs

◆ BlockScanT

◆ KeyValuePairT

◆ ReduceBySegmentOpT

◆ ScanTileStateT

◆ TilePrefixCallbackOpT

◆ ValueT

◆ WrappedFixupInputIteratorT

◆ WrappedPairsInputIteratorT

Member Enumeration Documentation

◆ anonymous enum

Constructor & Destructor Documentation

◆ AgentSegmentFixup()

Member Function Documentation

◆ ConsumeRange()

◆ ConsumeTile() [1/2]

◆ ConsumeTile() [2/2]

Field Documentation

◆ d_aggregates_out

◆ d_fixup_in

◆ d_pairs_in

◆ inequality_op

◆ reduction_op

◆ scan_op

◆ temp_storage