cub::TilePrefixCallbackOp#

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

Tile status interface for reduction by key, specialized for scan status and value types that can be combined into one machine word that can be read/written coherently in a single access.

Stateful block-scan prefix functor. Provides the the running prefix for the current tile by using the call-back warp to wait on on aggregates/prefixes from predecessor tiles to become available.

Template Parameters:

DelayConstructorT – Implementation detail, do not specify directly, requirements on the content of this type are subject to breaking change.

Public Types

using WarpReduceT = WarpReduce<T, (1 << (5))>#
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.

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#
struct TempStorage : public Uninitialized<_TempStorage>#