Enum Delivery 

pub enum Delivery {
    SyncPerUnit,
    AsyncPerUnit,
    AsyncBulk,
    Tma,
    Procedural,
}

Who moves an operand into a stage, and whether its bytes have landed when the fill returns.

Variants

SyncPerUnit

Every unit loads its own lines and stores them, through its registers: landed when the fill returns, and a line can be decoded on its way.

AsyncPerUnit

Every unit hands its own lines to the copy engine (cp.async), memory to memory: they land after the fill returns, as they lie, and something the slot waits on tracks them.

AsyncBulk

One unit hands the whole stage to the copy engine as one contiguous run (cp.async.bulk): it lands after the fill returns, counted in bytes on the slot’s barrier.

Tma

One unit hands the stage’s box to the TMA engine through a tensor map (cp.async.bulk.tensor): it lands after the fill returns, counted in bytes on the slot’s barrier.

Procedural

Every unit computes its values from their coordinates: nothing is read, and the stage holds them when the fill returns.

Implementations

impl Delivery

pub const TMA_STAGE_ALIGNMENT: usize = 1024

The byte alignment a stage a Tma box lands in starts on: the span a 128-byte swizzle repeats over, which the engine keys off the address, so a stage’s first row takes the first key as the fragment reading it does. A launch sizing its shared memory counts what aligning a stage to it can cost.

pub fn moves(self, stage: &Space) -> Result<(), Refusal>

Whether this delivery moves stage in one transfer, and why not where it cannot.

pub fn is_tma(&self) -> bool

Trait Implementations

impl Clone for Delivery

fn clone(&self) -> Self

Returns a duplicate of the value.

fn clone_from(&mut self, source: &Self)

Performs copy-assignment from source.

impl Copy for Delivery

impl Debug for Delivery

fn fmt(&self, f: &mut Formatter<'_>) -> Result

Formats the value using the given formatter.

impl Default for Delivery

fn default() -> Self

Returns the “default value” for a type.

impl Eq for Delivery

impl Hash for Delivery

fn hash<__H: Hasher>(&self, state: &mut __H)

Feeds this value into the given Hasher.

fn hash_slice<H>(data: &[Self], state: &mut H)

where H: Hasher, Self: Sized,

Feeds a slice of this type into the given Hasher.

impl PartialEq for Delivery

fn eq(&self, other: &Self) -> bool

Equality operator ==.

fn ne(&self, other: &Rhs) -> bool

Inequality operator !=.

impl StructuralPartialEq for Delivery