cub::BlockAdjacentDifference< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH > Class Template Reference

Detailed Description

template<typename T, int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>
class cub::BlockAdjacentDifference< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >

Definition at line 52 of file block_adjacent_difference.cuh.

Data Structures
struct	_TempStorage
	Shared memory storage layout type (last element from each thread's input) More...
struct	ApplyOp
	Specialization for when FlagOp has third index param. More...
struct	ApplyOp< FlagOp, false >
	Specialization for when FlagOp does not have a third index param. More...
struct	Iterate
	Templated unrolling of item comparison (inductive case) More...
struct	Iterate< MAX_ITERATIONS, MAX_ITERATIONS >
	Templated unrolling of item comparison (termination case) More...
struct	TempStorage
	\smemstorage{BlockDiscontinuity} More...

Public Member Functions
Collective constructors
__device__ __forceinline__	BlockAdjacentDifference ()
	Collective constructor using a private static allocation of shared memory as temporary storage.
__device__ __forceinline__	BlockAdjacentDifference (TempStorage &temp_storage)
	Collective constructor using the specified memory allocation as temporary storage.
Head flag operations
template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp >
__device__ __forceinline__ void	FlagHeads (FlagT(&head_flags)[ITEMS_PER_THREAD], T(&input)[ITEMS_PER_THREAD], T(&preds)[ITEMS_PER_THREAD], FlagOp flag_op)
template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp >
__device__ __forceinline__ void	FlagHeads (FlagT(&head_flags)[ITEMS_PER_THREAD], T(&input)[ITEMS_PER_THREAD], T(&preds)[ITEMS_PER_THREAD], FlagOp flag_op, T tile_predecessor_item)
template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp >
__device__ __forceinline__ void	FlagHeads (FlagT(&head_flags)[ITEMS_PER_THREAD], T(&input)[ITEMS_PER_THREAD], FlagOp flag_op)
template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp >
__device__ __forceinline__ void	FlagHeads (FlagT(&head_flags)[ITEMS_PER_THREAD], T(&input)[ITEMS_PER_THREAD], FlagOp flag_op, T tile_predecessor_item)
template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp >
__device__ __forceinline__ void	FlagTails (FlagT(&tail_flags)[ITEMS_PER_THREAD], T(&input)[ITEMS_PER_THREAD], FlagOp flag_op)
template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp >
__device__ __forceinline__ void	FlagTails (FlagT(&tail_flags)[ITEMS_PER_THREAD], T(&input)[ITEMS_PER_THREAD], FlagOp flag_op, T tile_successor_item)
template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp >
__device__ __forceinline__ void	FlagHeadsAndTails (FlagT(&head_flags)[ITEMS_PER_THREAD], FlagT(&tail_flags)[ITEMS_PER_THREAD], T(&input)[ITEMS_PER_THREAD], FlagOp flag_op)
template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp >
__device__ __forceinline__ void	FlagHeadsAndTails (FlagT(&head_flags)[ITEMS_PER_THREAD], FlagT(&tail_flags)[ITEMS_PER_THREAD], T tile_successor_item, T(&input)[ITEMS_PER_THREAD], FlagOp flag_op)
template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp >
__device__ __forceinline__ void	FlagHeadsAndTails (FlagT(&head_flags)[ITEMS_PER_THREAD], T tile_predecessor_item, FlagT(&tail_flags)[ITEMS_PER_THREAD], T(&input)[ITEMS_PER_THREAD], FlagOp flag_op)
template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp >
__device__ __forceinline__ void	FlagHeadsAndTails (FlagT(&head_flags)[ITEMS_PER_THREAD], T tile_predecessor_item, FlagT(&tail_flags)[ITEMS_PER_THREAD], T tile_successor_item, T(&input)[ITEMS_PER_THREAD], FlagOp flag_op)

Private Types
enum	{ BLOCK_THREADS = BLOCK_DIM_X * BLOCK_DIM_Y * BLOCK_DIM_Z }
	Constants. More...

Private Member Functions
__device__ __forceinline__ _TempStorage &	PrivateStorage ()
	Internal storage allocator.

Private Attributes
_TempStorage &	temp_storage
	Shared storage reference.
unsigned int	linear_tid
	Linear thread-id.

Member Enumeration Documentation

anonymous enum

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

anonymous enum

private

Constants.

Enumerator
BLOCK_THREADS	The thread block size in threads.

Definition at line 61 of file block_adjacent_difference.cuh.

Constructor & Destructor Documentation

BlockAdjacentDifference() [1/2]

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

__device__ __forceinline__ cub::BlockAdjacentDifference< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::BlockAdjacentDifference ( )

inline

Collective constructor using a private static allocation of shared memory as temporary storage.

Definition at line 215 of file block_adjacent_difference.cuh.

BlockAdjacentDifference() [2/2]

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

__device__ __forceinline__ cub::BlockAdjacentDifference< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::BlockAdjacentDifference ( TempStorage &  temp_storage )

inline

Collective constructor using the specified memory allocation as temporary storage.

Parameters

[in]

temp_storage

Reference to memory allocation having layout type TempStorage

Definition at line 225 of file block_adjacent_difference.cuh.

Member Function Documentation

FlagHeads() [1/4]

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp >

__device__ __forceinline__ void cub::BlockAdjacentDifference< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::FlagHeads ( FlagT(&)  head_flags[ITEMS_PER_THREAD], T(&)  input[ITEMS_PER_THREAD], FlagOp  flag_op  )

inline

Parameters

[out]	head_flags	Calling thread's discontinuity head_flags
[in]	input	Calling thread's input items
[in]	flag_op	Binary boolean flag predicate

Definition at line 306 of file block_adjacent_difference.cuh.

FlagHeads() [2/4]

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp >

__device__ __forceinline__ void cub::BlockAdjacentDifference< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::FlagHeads ( FlagT(&)  head_flags[ITEMS_PER_THREAD], T(&)  input[ITEMS_PER_THREAD], FlagOp  flag_op, T  tile_predecessor_item  )

inline

Parameters

[out]	head_flags	Calling thread's discontinuity head_flags
[in]	input	Calling thread's input items
[in]	flag_op	Binary boolean flag predicate
[in]	tile_predecessor_item	[thread0 only] Item with which to compare the first tile item (input0 from thread0).

Definition at line 320 of file block_adjacent_difference.cuh.

FlagHeads() [3/4]

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp >

__device__ __forceinline__ void cub::BlockAdjacentDifference< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::FlagHeads ( FlagT(&)  head_flags[ITEMS_PER_THREAD], T(&)  input[ITEMS_PER_THREAD], T(&)  preds[ITEMS_PER_THREAD], FlagOp  flag_op  )

inline

Parameters

[out]	head_flags	Calling thread's discontinuity head_flags
[in]	input	Calling thread's input items
[out]	preds	Calling thread's predecessor items
[in]	flag_op	Binary boolean flag predicate

Definition at line 246 of file block_adjacent_difference.cuh.

FlagHeads() [4/4]

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp >

__device__ __forceinline__ void cub::BlockAdjacentDifference< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::FlagHeads ( FlagT(&)  head_flags[ITEMS_PER_THREAD], T(&)  input[ITEMS_PER_THREAD], T(&)  preds[ITEMS_PER_THREAD], FlagOp  flag_op, T  tile_predecessor_item  )

inline

Parameters

[out]	head_flags	Calling thread's discontinuity head_flags
[in]	input	Calling thread's input items
[out]	preds	Calling thread's predecessor items
[in]	flag_op	Binary boolean flag predicate
[in]	tile_predecessor_item	[thread0 only] Item with which to compare the first tile item (input0 from thread0).

Definition at line 276 of file block_adjacent_difference.cuh.

FlagHeadsAndTails() [1/4]

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp >

__device__ __forceinline__ void cub::BlockAdjacentDifference< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::FlagHeadsAndTails ( FlagT(&)  head_flags[ITEMS_PER_THREAD], FlagT(&)  tail_flags[ITEMS_PER_THREAD], T  tile_successor_item, T(&)  input[ITEMS_PER_THREAD], FlagOp  flag_op  )

inline

Parameters

[out]	head_flags	Calling thread's discontinuity head_flags
[out]	tail_flags	Calling thread's discontinuity tail_flags
[in]	tile_successor_item	[threadBLOCK_THREADS-1 only] Item with which to compare the last tile item (inputITEMS_PER_THREAD-1 from threadBLOCK_THREADS-1).
[in]	input	Calling thread's input items
[in]	flag_op	Binary boolean flag predicate

Definition at line 446 of file block_adjacent_difference.cuh.

FlagHeadsAndTails() [2/4]

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp >

__device__ __forceinline__ void cub::BlockAdjacentDifference< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::FlagHeadsAndTails ( FlagT(&)  head_flags[ITEMS_PER_THREAD], FlagT(&)  tail_flags[ITEMS_PER_THREAD], T(&)  input[ITEMS_PER_THREAD], FlagOp  flag_op  )

inline

Parameters

[out]	head_flags	Calling thread's discontinuity head_flags
[out]	tail_flags	Calling thread's discontinuity tail_flags
[in]	input	Calling thread's input items
[in]	flag_op	Binary boolean flag predicate

Definition at line 395 of file block_adjacent_difference.cuh.

FlagHeadsAndTails() [3/4]

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp >

__device__ __forceinline__ void cub::BlockAdjacentDifference< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::FlagHeadsAndTails ( FlagT(&)  head_flags[ITEMS_PER_THREAD], T  tile_predecessor_item, FlagT(&)  tail_flags[ITEMS_PER_THREAD], T  tile_successor_item, T(&)  input[ITEMS_PER_THREAD], FlagOp  flag_op  )

inline

Parameters

[out]	head_flags	Calling thread's discontinuity head_flags
[in]	tile_predecessor_item	[thread0 only] Item with which to compare the first tile item (input0 from thread0).
[out]	tail_flags	Calling thread's discontinuity tail_flags
[in]	tile_successor_item	[threadBLOCK_THREADS-1 only] Item with which to compare the last tile item (inputITEMS_PER_THREAD-1 from threadBLOCK_THREADS-1).
[in]	input	Calling thread's input items
[in]	flag_op	Binary boolean flag predicate

Definition at line 545 of file block_adjacent_difference.cuh.

FlagHeadsAndTails() [4/4]

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp >

__device__ __forceinline__ void cub::BlockAdjacentDifference< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::FlagHeadsAndTails ( FlagT(&)  head_flags[ITEMS_PER_THREAD], T  tile_predecessor_item, FlagT(&)  tail_flags[ITEMS_PER_THREAD], T(&)  input[ITEMS_PER_THREAD], FlagOp  flag_op  )

inline

Parameters

[out]	head_flags	Calling thread's discontinuity head_flags
[in]	tile_predecessor_item	[thread0 only] Item with which to compare the first tile item (input0 from thread0).
[out]	tail_flags	Calling thread's discontinuity tail_flags
[in]	input	Calling thread's input items
[in]	flag_op	Binary boolean flag predicate

Definition at line 498 of file block_adjacent_difference.cuh.

FlagTails() [1/2]

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp >

__device__ __forceinline__ void cub::BlockAdjacentDifference< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::FlagTails ( FlagT(&)  tail_flags[ITEMS_PER_THREAD], T(&)  input[ITEMS_PER_THREAD], FlagOp  flag_op  )

inline

Parameters

[out]	tail_flags	Calling thread's discontinuity tail_flags
[in]	input	Calling thread's input items
[in]	flag_op	Binary boolean flag predicate

Definition at line 336 of file block_adjacent_difference.cuh.

FlagTails() [2/2]

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

template<int ITEMS_PER_THREAD, typename FlagT , typename FlagOp >

__device__ __forceinline__ void cub::BlockAdjacentDifference< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::FlagTails ( FlagT(&)  tail_flags[ITEMS_PER_THREAD], T(&)  input[ITEMS_PER_THREAD], FlagOp  flag_op, T  tile_successor_item  )

inline

Parameters

[out]	tail_flags	Calling thread's discontinuity tail_flags
[in]	input	Calling thread's input items
[in]	flag_op	Binary boolean flag predicate
[in]	tile_successor_item	[threadBLOCK_THREADS-1 only] Item with which to compare the last tile item (inputITEMS_PER_THREAD-1 from threadBLOCK_THREADS-1).

Definition at line 364 of file block_adjacent_difference.cuh.

PrivateStorage()

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

__device__ __forceinline__ _TempStorage & cub::BlockAdjacentDifference< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::PrivateStorage ( )

inlineprivate

Internal storage allocator.

Definition at line 81 of file block_adjacent_difference.cuh.

Field Documentation

linear_tid

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

unsigned int cub::BlockAdjacentDifference< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::linear_tid

private

Linear thread-id.

Definition at line 198 of file block_adjacent_difference.cuh.

temp_storage

template<typename T , int BLOCK_DIM_X, int BLOCK_DIM_Y = 1, int BLOCK_DIM_Z = 1, int PTX_ARCH = CUB_PTX_ARCH>

_TempStorage& cub::BlockAdjacentDifference< T, BLOCK_DIM_X, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH >::temp_storage

private

Shared storage reference.

Definition at line 195 of file block_adjacent_difference.cuh.

---

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

• openfpm_data/src/util/cuda/cub_old/block/block_adjacent_difference.cuh