cub::TilePrefixCallbackOp#

template<typename T, typename ScanOpT, typename ScanTileStateT, typename DelayConstructorT = detail::default_delay_constructor_t<T>, bool StableReductionOrder = false>
struct TilePrefixCallbackOp#

Public Types

using WarpReduceT = WarpReduce<T, (1 << (5))>#
using TempStorage = Uninitialized<_TempStorage>#
using StatusWord = typename ScanTileStateT::StatusWord#

Public Functions

inline TilePrefixCallbackOp(
ScanTileStateT &tile_status,
TempStorage &temp_storage,
ScanOpT scan_op,
int tile_idx
)#
inline TilePrefixCallbackOp(
ScanTileStateT &tile_status,
TempStorage &temp_storage,
ScanOpT scan_op
)#
template<class DelayT = detail::default_delay_t<T>>
inline void ProcessWindow(
int predecessor_idx,
StatusWord &predecessor_status,
T &window_aggregate,
DelayT delay = {}
)#

Block until all predecessors within the warp-wide window have non-invalid status.

Parameters:

  • predecessor_idx – Preceding tile index to inspect

  • predecessor_status – [out] Preceding tile status

  • window_aggregate – [out] Relevant partial reduction from this window of preceding tiles

inline T operator()(T block_aggregate)#
inline T GetExclusivePrefix()#
inline T GetInclusivePrefix()#
inline T GetBlockAggregate()#
inline int GetTileIdx() const#

Public Members

_TempStorage &temp_storage#

Reference to a warp-reduction instance.

ScanTileStateT &tile_status#

Interface to tile status (non-anchor tiles stay at PARTIAL on the deterministic path)

ScanOpT scan_op#

Binary scan operator.

int tile_idx#

The current tile index.

T exclusive_prefix#

Exclusive prefix for the tile.

T inclusive_prefix#

Inclusive prefix for the tile.

struct _TempStorage#

Public Members

WarpReduceT::TempStorage warp_reduce#
T exclusive_prefix#
T inclusive_prefix#
T block_aggregate#