For the complete documentation index, see llms.txt. Markdown versions of all pages are available by appending .md to any URL (e.g. /get-started.md).
Mojo trait
BarrierHandler
Handles barrier lifecycle for different transfer mechanisms.
Separates barrier management from tile loading:
- prepare_stage: Called once before loading tiles for a stage.
- complete_stage: Called once after all tiles for a stage are loaded.
TMA: prepare sets expected bytes, complete is noop (hardware signals). cp.async: prepare is noop, complete commits copies and signals arrival.
Implemented traits
AnyType,
Copyable,
Deinitable,
ImplicitlyCopyable,
Movable,
RegisterPassable,
TrivialRegisterPassable
Required methods
prepare_stage
def prepare_stage(self, mem_barrier: Pointer[SharedMemBarrier, address_space=AddressSpace.SHARED])
Prepare barrier for incoming transfers.
For TMA: sets expected transaction bytes. For cp.async: noop.
Args:
- mem_barrier (
Pointer[SharedMemBarrier, address_space=AddressSpace.SHARED]): The stage's memory barrier.
complete_stage
def complete_stage(self, mem_barrier: Pointer[SharedMemBarrier, address_space=AddressSpace.SHARED])
Signal that all transfers for this stage are done.
For TMA: noop (hardware auto-signals). For cp.async: commits pending copies and signals thread arrival.
Args:
- mem_barrier (
Pointer[SharedMemBarrier, address_space=AddressSpace.SHARED]): The stage's memory barrier.