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.

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

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

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

Constructor & Destructor Documentation

◆ AgentSegmentFixup()

Member Function Documentation

◆ ConsumeRange()

◆ ConsumeTile() [1/2]

◆ ConsumeTile() [2/2]