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,
ImplicitlyCopyable,
ImplicitlyDestructible,
Movable,
RegisterPassable,
TrivialRegisterPassable
comptime members
__copyinit__is_trivial
comptime __copyinit__is_trivial
A flag (often compiler generated) to indicate whether the implementation of __copyinit__ is trivial.
The implementation of __copyinit__ is considered to be trivial if:
- The struct has a compiler-generated trivial
__copyinit__and all its fields have a trivial__copyinit__method.
In practice, it means the value can be copied by copying the bits from one location to another without side effects.
__del__is_trivial
comptime __del__is_trivial
A flag (often compiler generated) to indicate whether the implementation of __del__ is trivial.
The implementation of __del__ is considered to be trivial if:
- The struct has a compiler-generated trivial destructor and all its fields
have a trivial
__del__method.
In practice, it means that the __del__ can be considered as no-op.
__moveinit__is_trivial
comptime __moveinit__is_trivial
A flag (often compiler generated) to indicate whether the implementation of __moveinit__ is trivial.
The implementation of __moveinit__ is considered to be trivial if:
- The struct has a compiler-generated
__moveinit__and all its fields have a trivial__moveinit__method.
In practice, it means the value can be moved by moving the bits from one location to another without side effects.
Required methods
__copyinit__
__copyinit__(out self: _Self, existing: _Self, /)
Create a new instance of the value by copying an existing one.
Args:
- existing (
_Self): The value to copy.
Returns:
_Self
__moveinit__
__moveinit__(out self: _Self, deinit existing: _Self, /)
Create a new instance of the value by moving the value of another.
Args:
- existing (
_Self): The value to move.
Returns:
_Self
prepare_stage
prepare_stage(self: _Self, mem_barrier: LegacyUnsafePointer[SharedMemBarrier, address_space=AddressSpace.SHARED])
Prepare barrier for incoming transfers.
For TMA: sets expected transaction bytes. For cp.async: noop.
Args:
- mem_barrier (
LegacyUnsafePointer): The stage's memory barrier.
complete_stage
complete_stage(self: _Self, mem_barrier: LegacyUnsafePointer[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 (
LegacyUnsafePointer): The stage's memory barrier.
Provided methods
copy
copy(self: _Self) -> _Self
Explicitly construct a copy of self.
Returns:
_Self: A copy of this value.
Was this page helpful?
Thank you! We'll create more content like this.
Thank you for helping us improve!