AgentRle implements a stateful abstraction of CUDA thread blocks for participating in device-wide run-length-encode. More...

Detailed Description

template<typename AgentRlePolicyT, typename InputIteratorT, typename OffsetsOutputIteratorT, typename LengthsOutputIteratorT, typename EqualityOpT, typename OffsetT>
struct cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >

AgentRle implements a stateful abstraction of CUDA thread blocks for participating in device-wide run-length-encode.

< Signed integer type for global offsets

Definition at line 102 of file agent_rle.cuh.

Data Structures
struct	_TempStorage

struct	OobInequalityOp

struct	TempStorage

Public Types
enum	{ WARP_THREADS = CUB_WARP_THREADS(PTX_ARCH) , BLOCK_THREADS = AgentRlePolicyT::BLOCK_THREADS , ITEMS_PER_THREAD = AgentRlePolicyT::ITEMS_PER_THREAD , WARP_ITEMS = WARP_THREADS * ITEMS_PER_THREAD , TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD , WARPS = (BLOCK_THREADS + WARP_THREADS - 1) / WARP_THREADS , SYNC_AFTER_LOAD = (AgentRlePolicyT::LOAD_ALGORITHM != BLOCK_LOAD_DIRECT) , STORE_WARP_TIME_SLICING = AgentRlePolicyT::STORE_WARP_TIME_SLICING , ACTIVE_EXCHANGE_WARPS = (STORE_WARP_TIME_SLICING) ? 1 : WARPS }

typedef std::iterator_traits< InputIteratorT >::value_type	T
	The input value type.

typedef If<(Equals< typenamestd::iterator_traits< LengthsOutputIteratorT >::value_type, void >::VALUE), OffsetT, typenamestd::iterator_traits< LengthsOutputIteratorT >::value_type >::Type	LengthT
	The lengths output value type.

typedef KeyValuePair< OffsetT, LengthT >	LengthOffsetPair
	Tuple type for scanning (pairs run-length and run-index)

typedef ReduceByKeyScanTileState< LengthT, OffsetT >	ScanTileStateT
	Tile status descriptor interface type.

typedef If< IsPointer< InputIteratorT >::VALUE, CacheModifiedInputIterator< AgentRlePolicyT::LOAD_MODIFIER, T, OffsetT >, InputIteratorT >::Type	WrappedInputIteratorT

typedef BlockLoad< T, AgentRlePolicyT::BLOCK_THREADS, AgentRlePolicyT::ITEMS_PER_THREAD, AgentRlePolicyT::LOAD_ALGORITHM >	BlockLoadT

typedef BlockDiscontinuity< T, BLOCK_THREADS >	BlockDiscontinuityT

typedef WarpScan< LengthOffsetPair >	WarpScanPairs

typedef ReduceBySegmentOp< cub::Sum >	ReduceBySegmentOpT

typedef TilePrefixCallbackOp< LengthOffsetPair, ReduceBySegmentOpT, ScanTileStateT >	TilePrefixCallbackOpT

typedef WarpExchange< LengthOffsetPair, ITEMS_PER_THREAD >	WarpExchangePairs

typedef If< STORE_WARP_TIME_SLICING, typenameWarpExchangePairs::TempStorage, NullType >::Type	WarpExchangePairsStorage

typedef WarpExchange< OffsetT, ITEMS_PER_THREAD >	WarpExchangeOffsets

typedef WarpExchange< LengthT, ITEMS_PER_THREAD >	WarpExchangeLengths

typedef LengthOffsetPair	WarpAggregates[WARPS]

Public Member Functions
__device__ __forceinline__	AgentRle (TempStorage &temp_storage, InputIteratorT d_in, OffsetsOutputIteratorT d_offsets_out, LengthsOutputIteratorT d_lengths_out, EqualityOpT equality_op, OffsetT num_items)

template<bool FIRST_TILE, bool LAST_TILE>
__device__ __forceinline__ void	InitializeSelections (OffsetT tile_offset, OffsetT num_remaining, T(&items)[ITEMS_PER_THREAD], LengthOffsetPair(&lengths_and_num_runs)[ITEMS_PER_THREAD])

__device__ __forceinline__ void	WarpScanAllocations (LengthOffsetPair &tile_aggregate, LengthOffsetPair &warp_aggregate, LengthOffsetPair &warp_exclusive_in_tile, LengthOffsetPair &thread_exclusive_in_warp, LengthOffsetPair(&lengths_and_num_runs)[ITEMS_PER_THREAD])

template<bool FIRST_TILE>
__device__ __forceinline__ void	ScatterTwoPhase (OffsetT tile_num_runs_exclusive_in_global, OffsetT warp_num_runs_aggregate, OffsetT warp_num_runs_exclusive_in_tile, OffsetT(&thread_num_runs_exclusive_in_warp)[ITEMS_PER_THREAD], LengthOffsetPair(&lengths_and_offsets)[ITEMS_PER_THREAD], Int2Type< true > is_warp_time_slice)

template<bool FIRST_TILE>
__device__ __forceinline__ void	ScatterTwoPhase (OffsetT tile_num_runs_exclusive_in_global, OffsetT warp_num_runs_aggregate, OffsetT warp_num_runs_exclusive_in_tile, OffsetT(&thread_num_runs_exclusive_in_warp)[ITEMS_PER_THREAD], LengthOffsetPair(&lengths_and_offsets)[ITEMS_PER_THREAD], Int2Type< false > is_warp_time_slice)

template<bool FIRST_TILE>
__device__ __forceinline__ void	ScatterDirect (OffsetT tile_num_runs_exclusive_in_global, OffsetT warp_num_runs_aggregate, OffsetT warp_num_runs_exclusive_in_tile, OffsetT(&thread_num_runs_exclusive_in_warp)[ITEMS_PER_THREAD], LengthOffsetPair(&lengths_and_offsets)[ITEMS_PER_THREAD])

template<bool FIRST_TILE>
__device__ __forceinline__ void	Scatter (OffsetT tile_num_runs_aggregate, OffsetT tile_num_runs_exclusive_in_global, OffsetT warp_num_runs_aggregate, OffsetT warp_num_runs_exclusive_in_tile, OffsetT(&thread_num_runs_exclusive_in_warp)[ITEMS_PER_THREAD], LengthOffsetPair(&lengths_and_offsets)[ITEMS_PER_THREAD])

template<bool LAST_TILE>
__device__ __forceinline__ LengthOffsetPair	ConsumeTile (OffsetT num_items, OffsetT num_remaining, int tile_idx, OffsetT tile_offset, ScanTileStateT &tile_status)

template<typename NumRunsIteratorT >
__device__ __forceinline__ void	ConsumeRange (int num_tiles, ScanTileStateT &tile_status, NumRunsIteratorT d_num_runs_out)
	< Output iterator type for recording number of items selected

Data Fields
_TempStorage &	temp_storage
	Reference to temp_storage.

WrappedInputIteratorT	d_in
	Pointer to input sequence of data items.

OffsetsOutputIteratorT	d_offsets_out
	Input run offsets.

LengthsOutputIteratorT	d_lengths_out
	Output run lengths.

EqualityOpT	equality_op
	T equality operator.

ReduceBySegmentOpT	scan_op
	Reduce-length-by-flag scan operator.

OffsetT	num_items
	Total number of input items.

Member Typedef Documentation

◆ BlockDiscontinuityT

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

typedef BlockDiscontinuity<T, BLOCK_THREADS> cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::BlockDiscontinuityT

Definition at line 186 of file agent_rle.cuh.

◆ BlockLoadT

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

typedef BlockLoad< T, AgentRlePolicyT::BLOCK_THREADS, AgentRlePolicyT::ITEMS_PER_THREAD, AgentRlePolicyT::LOAD_ALGORITHM> cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::BlockLoadT

Definition at line 183 of file agent_rle.cuh.

◆ LengthOffsetPair

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

typedef KeyValuePair<OffsetT, LengthT> cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::LengthOffsetPair

Tuple type for scanning (pairs run-length and run-index)

Definition at line 117 of file agent_rle.cuh.

◆ LengthT

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

typedef If<(Equals<typenamestd::iterator_traits<LengthsOutputIteratorT>::value_type,void>::VALUE),OffsetT,typenamestd::iterator_traits<LengthsOutputIteratorT>::value_type>::Type cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::LengthT

The lengths output value type.

Definition at line 114 of file agent_rle.cuh.

◆ ReduceBySegmentOpT

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

typedef ReduceBySegmentOp<cub::Sum> cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::ReduceBySegmentOpT

Definition at line 192 of file agent_rle.cuh.

◆ ScanTileStateT

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

typedef ReduceByKeyScanTileState<LengthT, OffsetT> cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::ScanTileStateT

Tile status descriptor interface type.

Definition at line 120 of file agent_rle.cuh.

◆ T

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

typedef std::iterator_traits<InputIteratorT>::value_type cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::T

The input value type.

Definition at line 109 of file agent_rle.cuh.

◆ TilePrefixCallbackOpT

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

typedef TilePrefixCallbackOp< LengthOffsetPair, ReduceBySegmentOpT, ScanTileStateT> cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::TilePrefixCallbackOpT

Definition at line 199 of file agent_rle.cuh.

◆ WarpAggregates

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

typedef LengthOffsetPair cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::WarpAggregates[WARPS]

Definition at line 209 of file agent_rle.cuh.

◆ WarpExchangeLengths

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

typedef WarpExchange<LengthT, ITEMS_PER_THREAD> cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::WarpExchangeLengths

Definition at line 207 of file agent_rle.cuh.

◆ WarpExchangeOffsets

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

typedef WarpExchange<OffsetT, ITEMS_PER_THREAD> cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::WarpExchangeOffsets

Definition at line 206 of file agent_rle.cuh.

◆ WarpExchangePairs

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

typedef WarpExchange<LengthOffsetPair, ITEMS_PER_THREAD> cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::WarpExchangePairs

Definition at line 202 of file agent_rle.cuh.

◆ WarpExchangePairsStorage

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

typedef If<STORE_WARP_TIME_SLICING,typenameWarpExchangePairs::TempStorage,NullType>::Type cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::WarpExchangePairsStorage

Definition at line 204 of file agent_rle.cuh.

◆ WarpScanPairs

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

typedef WarpScan<LengthOffsetPair> cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::WarpScanPairs

Definition at line 189 of file agent_rle.cuh.

◆ WrappedInputIteratorT

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

typedef If<IsPointer<InputIteratorT>::VALUE,CacheModifiedInputIterator<AgentRlePolicyT::LOAD_MODIFIER,T,OffsetT>,InputIteratorT>::Type cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::WrappedInputIteratorT

Definition at line 175 of file agent_rle.cuh.

Member Enumeration Documentation

◆ anonymous enum

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

anonymous enum

Enumerator
SYNC_AFTER_LOAD	Whether or not to sync after loading data.
STORE_WARP_TIME_SLICING	Whether or not only one warp's worth of shared memory should be allocated and time-sliced among block-warps during any store-related data transpositions (versus each warp having its own storage)

Definition at line 123 of file agent_rle.cuh.

Constructor & Destructor Documentation

◆ AgentRle()

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

__device__ __forceinline__ cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::AgentRle	(	TempStorage &	temp_storage,
		InputIteratorT	d_in,
		OffsetsOutputIteratorT	d_offsets_out,
		LengthsOutputIteratorT	d_lengths_out,
		EqualityOpT	equality_op,
		OffsetT	num_items
	)

inline

Parameters

[in]	temp_storage	Reference to temp_storage
[in]	d_in	Pointer to input sequence of data items
[out]	d_offsets_out	Pointer to output sequence of run offsets
[out]	d_lengths_out	Pointer to output sequence of run lengths
[in]	equality_op	T equality operator
[in]	num_items	Total number of input items

Definition at line 270 of file agent_rle.cuh.

Member Function Documentation

◆ ConsumeRange()

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

template<typename NumRunsIteratorT >

__device__ __forceinline__ void cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::ConsumeRange	(	int	num_tiles,
		ScanTileStateT &	tile_status,
		NumRunsIteratorT	d_num_runs_out
	)

inline

< Output iterator type for recording number of items selected

Scan tiles of items as part of a dynamic chained scan

Parameters

num_tiles	Total number of input tiles
tile_status	Global list of tile status
d_num_runs_out	Output pointer for total number of runs identified

Definition at line 801 of file agent_rle.cuh.

◆ ConsumeTile()

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

template<bool LAST_TILE>

__device__ __forceinline__ LengthOffsetPair cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::ConsumeTile	(	OffsetT	num_items,
		OffsetT	num_remaining,
		int	tile_idx,
		OffsetT	tile_offset,
		ScanTileStateT &	tile_status
	)

inline

Process a tile of input (dynamic chained scan)

Parameters

num_items	Total number of global input items
num_remaining	Number of global input items remaining (including this tile)
tile_idx	Tile index
tile_offset	Tile offset
tile_status	Global list of tile status

Definition at line 617 of file agent_rle.cuh.

◆ InitializeSelections()

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

template<bool FIRST_TILE, bool LAST_TILE>

__device__ __forceinline__ void cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::InitializeSelections	(	OffsetT	tile_offset,
		OffsetT	num_remaining,
		T(&)	items[ITEMS_PER_THREAD],
		LengthOffsetPair(&)	lengths_and_num_runs[ITEMS_PER_THREAD]
	)

inline

Definition at line 293 of file agent_rle.cuh.

◆ Scatter()

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

template<bool FIRST_TILE>

__device__ __forceinline__ void cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::Scatter	(	OffsetT	tile_num_runs_aggregate,
		OffsetT	tile_num_runs_exclusive_in_global,
		OffsetT	warp_num_runs_aggregate,
		OffsetT	warp_num_runs_exclusive_in_tile,
		OffsetT(&)	thread_num_runs_exclusive_in_warp[ITEMS_PER_THREAD],
		LengthOffsetPair(&)	lengths_and_offsets[ITEMS_PER_THREAD]
	)

inline

Scatter

Definition at line 572 of file agent_rle.cuh.

◆ ScatterDirect()

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

template<bool FIRST_TILE>

__device__ __forceinline__ void cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::ScatterDirect	(	OffsetT	tile_num_runs_exclusive_in_global,
		OffsetT	warp_num_runs_aggregate,
		OffsetT	warp_num_runs_exclusive_in_tile,
		OffsetT(&)	thread_num_runs_exclusive_in_warp[ITEMS_PER_THREAD],
		LengthOffsetPair(&)	lengths_and_offsets[ITEMS_PER_THREAD]
	)

inline

Direct scatter

Definition at line 538 of file agent_rle.cuh.

◆ ScatterTwoPhase() [1/2]

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

template<bool FIRST_TILE>

__device__ __forceinline__ void cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::ScatterTwoPhase	(	OffsetT	tile_num_runs_exclusive_in_global,
		OffsetT	warp_num_runs_aggregate,
		OffsetT	warp_num_runs_exclusive_in_tile,
		OffsetT(&)	thread_num_runs_exclusive_in_warp[ITEMS_PER_THREAD],
		LengthOffsetPair(&)	lengths_and_offsets[ITEMS_PER_THREAD],
		Int2Type< false >	is_warp_time_slice
	)

inline

Two-phase scatter

Definition at line 480 of file agent_rle.cuh.

◆ ScatterTwoPhase() [2/2]

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

template<bool FIRST_TILE>

__device__ __forceinline__ void cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::ScatterTwoPhase	(	OffsetT	tile_num_runs_exclusive_in_global,
		OffsetT	warp_num_runs_aggregate,
		OffsetT	warp_num_runs_exclusive_in_tile,
		OffsetT(&)	thread_num_runs_exclusive_in_warp[ITEMS_PER_THREAD],
		LengthOffsetPair(&)	lengths_and_offsets[ITEMS_PER_THREAD],
		Int2Type< true >	is_warp_time_slice
	)

inline

Two-phase scatter, specialized for warp time-slicing

Definition at line 421 of file agent_rle.cuh.

◆ WarpScanAllocations()

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

__device__ __forceinline__ void cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::WarpScanAllocations	(	LengthOffsetPair &	tile_aggregate,
		LengthOffsetPair &	warp_aggregate,
		LengthOffsetPair &	warp_exclusive_in_tile,
		LengthOffsetPair &	thread_exclusive_in_warp,
		LengthOffsetPair(&)	lengths_and_num_runs[ITEMS_PER_THREAD]
	)

inline

Scan of allocations

Definition at line 367 of file agent_rle.cuh.

Field Documentation

◆ d_in

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

WrappedInputIteratorT cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::d_in

Pointer to input sequence of data items.

Definition at line 255 of file agent_rle.cuh.

◆ d_lengths_out

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

LengthsOutputIteratorT cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::d_lengths_out

Output run lengths.

Definition at line 257 of file agent_rle.cuh.

◆ d_offsets_out

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

OffsetsOutputIteratorT cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::d_offsets_out

Input run offsets.

Definition at line 256 of file agent_rle.cuh.

◆ equality_op

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

EqualityOpT cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::equality_op

T equality operator.

Definition at line 259 of file agent_rle.cuh.

◆ num_items

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

OffsetT cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::num_items

Total number of input items.

Definition at line 261 of file agent_rle.cuh.

◆ scan_op

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

ReduceBySegmentOpT cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::scan_op

Reduce-length-by-flag scan operator.

Definition at line 260 of file agent_rle.cuh.

◆ temp_storage

template<typename AgentRlePolicyT , typename InputIteratorT , typename OffsetsOutputIteratorT , typename LengthsOutputIteratorT , typename EqualityOpT , typename OffsetT >

_TempStorage& cub::AgentRle< AgentRlePolicyT, InputIteratorT, OffsetsOutputIteratorT, LengthsOutputIteratorT, EqualityOpT, OffsetT >::temp_storage

Reference to temp_storage.

Definition at line 253 of file agent_rle.cuh.

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

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

Detailed Description

Data Structures

Public Types

Public Member Functions

Data Fields

Member Typedef Documentation

◆ BlockDiscontinuityT

◆ BlockLoadT

◆ LengthOffsetPair

◆ LengthT

◆ ReduceBySegmentOpT

◆ ScanTileStateT

◆ T

◆ TilePrefixCallbackOpT

◆ WarpAggregates

◆ WarpExchangeLengths

◆ WarpExchangeOffsets

◆ WarpExchangePairs

◆ WarpExchangePairsStorage

◆ WarpScanPairs

◆ WrappedInputIteratorT

Member Enumeration Documentation

◆ anonymous enum

Constructor & Destructor Documentation

◆ AgentRle()

Member Function Documentation

◆ ConsumeRange()

◆ ConsumeTile()

◆ InitializeSelections()

◆ Scatter()

◆ ScatterDirect()

◆ ScatterTwoPhase() [1/2]

◆ ScatterTwoPhase() [2/2]

◆ WarpScanAllocations()

Field Documentation

◆ d_in

◆ d_lengths_out

◆ d_offsets_out

◆ equality_op

◆ num_items

◆ scan_op

◆ temp_storage