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
__copy_ctor_is_trivial
comptime __copy_ctor_is_trivial
A flag (often compiler generated) to indicate whether the implementation of the copy constructor is trivial.
A copy constructor is considered to be trivial if:
- The struct has a compiler-generated trivial copy constructor because all its fields have trivial copy constructors.
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.
__move_ctor_is_trivial
comptime __move_ctor_is_trivial
A flag (often compiler generated) to indicate whether the implementation of move constructor is trivial.
The implementation of a move constructor is considered to be trivial if:
- The struct has a compiler-generated trivial move constructor because all its fields have trivial move constructors.
In practice, it means the value can be moved by moving the bits from one location to another without side effects.
Required methods
__init__
__init__(out self: _Self, *, copy: _Self)
Create a new instance of the value by copying an existing one.
Args:
- copy (
_Self): The value to copy.
Returns:
_Self
__init__(out self: _Self, *, deinit take: _Self)
Create a new instance of the value by moving the value of another.
Args:
- take (
_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, a convenience method for Self(copy=self) when the type is inconvenient to write out.
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!