cub::ScanTileState< T, false >#

template<typename T>
struct ScanTileState<T, false>#

Tile status interface 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.

Tile status interface specialized for scan status and value types that cannot be combined into one machine word.

Public Types

enum [anonymous]#

Values:

enumerator TILE_STATUS_PADDING = detail::warp_threads#
using StatusValueT = T#
using StatusWord = unsigned int#

Public Functions

inline ScanTileState()#

Constructor.

inline cudaError_t Init(
int num_tiles,
void *d_temp_storage,
size_t temp_storage_bytes,
)#

Initializer.

Parameters:

  • num_tiles – [in] Number of tiles

  • d_temp_storage – [in] Device-accessible allocation of temporary storage. When nullptr, the required allocation size is written to temp_storage_bytes and no work is done.

  • temp_storage_bytes – [in] Size in bytes of d_temp_storage allocation Initializer

inline void InitializeStatus(int num_tiles)#

Initialize (from device)

template<MemoryOrder Order = MemoryOrder::relaxed>
inline void SetInclusive(
int tile_idx,
T tile_inclusive,
)#

Update the specified tile’s inclusive value and corresponding status.

template<MemoryOrder Order = MemoryOrder::relaxed>
inline void SetPartial(
int tile_idx,
T tile_partial,
)#

Update the specified tile’s partial value and corresponding status.

template<class DelayT = detail::default_no_delay_t, MemoryOrder Order = MemoryOrder::relaxed>
inline void WaitForValid(
int tile_idx,
StatusWord &status,
T &value,
DelayT delay = {},
)#

Wait for the corresponding tile to become non-invalid.

inline T LoadValid(int tile_idx)#

Loads and returns the tile’s value.

The returned value is undefined if either (a) the tile’s status is invalid or (b) there is no memory fence between reading a non-invalid status and the call to LoadValid.

Public Members

StatusWord *d_tile_status#
T *d_tile_partial#
T *d_tile_inclusive#

Public Static Functions

static inline constexpr cudaError_t AllocationSize(
int num_tiles,
size_t &temp_storage_bytes,
)#

Compute device memory needed for tile status.

Parameters:

  • num_tiles – [in] Number of tiles

  • temp_storage_bytes – [out] Size in bytes of d_temp_storage allocation

Public Static Attributes

static constexpr size_t description_bytes_per_tile = sizeof(StatusWord)#
static constexpr size_t payload_bytes_per_tile = sizeof(Uninitialized<T>)#