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< typename std::iterator_traits< LengthsOutputIteratorT >::value_type, void >::VALUE), OffsetT, typename std::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, typename WarpExchangePairs::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 More...

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

◆ 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< true >	is_warp_time_slice
	)

inline

Two-phase scatter, specialized for warp time-slicing

Definition at line 421 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< false >	is_warp_time_slice
	)

inline

Two-phase scatter

Definition at line 480 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.

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

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

Detailed Description

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

Data Structures

Public Types

Public Member Functions

Data Fields

Member Enumeration Documentation

◆ anonymous enum

Constructor & Destructor Documentation

◆ AgentRle()

Member Function Documentation

◆ ConsumeRange()

◆ ConsumeTile()

◆ Scatter()

◆ ScatterDirect()

◆ ScatterTwoPhase() [1/2]

◆ ScatterTwoPhase() [2/2]

◆ WarpScanAllocations()

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