cub::RleLookaheadPolicy

struct RleLookaheadPolicy

The tuning policy for the lookahead implementation of DeviceRunLengthEncode::Encode.

Public Functions

inline constexpr int warp_tile_size() const noexcept

inline constexpr int tile_size() const noexcept

inline constexpr int decode_items_per_thread() const noexcept

inline constexpr int slot_pad(int key_size) const noexcept

inline constexpr int slot_stride(

int key_size,

int key_align

) const noexcept

inline constexpr ::cuda::std::size_t dyn_smem_bytes(

int key_size,

int key_align

) const noexcept

inline constexpr int floor_key_ring_stages() const noexcept

inline constexpr int floor_pos_ring_stages() const noexcept

inline constexpr ::cuda::std::size_t floor_dyn_smem_bytes(

) const noexcept

inline friend ::std::ostream &operator<<(

::std::ostream &os,

const RleLookaheadPolicy &p

)

Public Members

int items_per_thread

Number of items each lane of a compute warp processes; a warp tile is warp_threads * items_per_thread items.

int compute_warps

Number of compute warps; each processes one warp tile per pipeline generation.

int key_ring_stages

Depth of the key staging ring: how many pipeline generations can be in flight.

int pos_ring_stages

Depth of the run-positions ring; 2 * pos_ring_stages >= key_ring_stages must hold.

int poll_items_per_thread

Number of tile-state loads each poll-warp lane keeps in flight.

int dense_poll_items_per_thread

Loads per lane for the smaller poll window used in dense mode; the window is warp_threads * dense_poll_items_per_thread tile states.

int dense_mode_runs_per_tile

Average runs per tile in a folded window above which the poll switches to the smaller dense-mode window.

int flag_staging_threshold

Runs per warp tile below which the compute warp stages raw head flags and the store warp decodes positions itself, instead of staging precomputed positions.

Public Static Attributes

static constexpr ::cuda::std::size_t static_smem_budget = 8 * 1024

the unstaged floor configuration keeps at most this many key generations in flight

static constexpr int floor_key_ring_cap = 4

one pos-ring stage may cover at most this many key generations: the parity bound pos_ring_stages * max_key_stages_per_pos_stage >= key_ring_stages must hold

static constexpr int max_key_stages_per_pos_stage = 2

Friends

inline friend constexpr bool operator==(

const RleLookaheadPolicy &lhs,

const RleLookaheadPolicy &rhs

) noexcept

inline friend constexpr bool operator!=(

const RleLookaheadPolicy &lhs,

const RleLookaheadPolicy &rhs

) noexcept