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 struct
NvidiaMbarBackend
struct NvidiaMbarBackend[num_stages: Int]
PipelineBackend using NVIDIA mbarrier objects.
This is the default backend and reproduces the historical, hardcoded
mbarrier behavior exactly. The ring's 2 * num_stages SharedMemBarrier
objects are laid out as full[0..num_stages) followed by
empty[0..num_stages); full points at the first, empty at the second.
Parameters
- num_stages (
Int): The number of pipeline stages (ring depth) this backend instance is configured for.
Fields
- full (
MbarPtr): - empty (
MbarPtr):
Implemented traits
AnyType,
Copyable,
Deinitable,
ImplicitlyCopyable,
Movable,
PipelineBackend,
RegisterPassable,
TrivialRegisterPassable
comptime members
BarrierStorage
comptime BarrierStorage = SharedMemBarrier
Handle
comptime Handle = MbarPtr
Methods
__init__
def __init__[passed_num_stages: Int](ptr: Pointer[SharedMemBarrier, MutUntrackedOrigin, address_space=AddressSpace.SHARED]) -> Self
Construct from the base pointer of the backing barrier array.
Parameters:
- passed_num_stages (
Int): The number of pipeline stages (ring depth). Must matchSelf.num_stages.
Args:
- ptr (
Pointer[SharedMemBarrier, MutUntrackedOrigin, address_space=AddressSpace.SHARED]): Pointer to the first of2 * num_stagesbarriers.
storage_elems
init_barriers
def init_barriers[passed_num_stages: Int](self, producer_arrive_count: Int32, consumer_arrive_count: Int32)
wait_full
def wait_full[ticks: Optional[UInt32] = None](self, stage: UInt32, phase: UInt32)
wait_empty
def wait_empty[ticks: Optional[UInt32] = None](self, stage: UInt32, phase: UInt32)
try_full
try_empty
arrive_full
def arrive_full(self, stage: UInt32)
arrive_empty
def arrive_empty(self, stage: UInt32)
full_handle
def full_handle(self, stage: UInt32) -> Self.Handle
Returns:
Self.Handle
empty_handle
def empty_handle(self, stage: UInt32) -> Self.Handle
Returns:
Self.Handle