Mojo struct
RingBuffer
@register_passable(trivial)
struct RingBuffer[num_pipeline_stages: Int, num_consumers: Int, cluster_size: Int, use_async_copy: Bool = False]
Ring buffer for managing pipeline synchronization between producers and consumers.
This struct encapsulates the synchronization logic for a multi-stage pipeline with one producer and multiple consumers, supporting both single-block and multi-cluster configurations.
Fields
- full_mbar (
UnsafePointer[SharedMemBarrier, address_space=AddressSpace(3)]): - empty_mbar (
UnsafePointer[SharedMemBarrier, address_space=AddressSpace(3)]): - read_state (
PipelineState[num_pipeline_stages]): - write_state (
PipelineState[num_pipeline_stages]): - warp_group_thread_idx (
UInt):
Implemented traits
AnyType,
Copyable,
ImplicitlyCopyable,
Movable,
UnknownDestructibility
Aliases
__copyinit__is_trivial
alias __copyinit__is_trivial = True
__del__is_trivial
alias __del__is_trivial = True
__moveinit__is_trivial
alias __moveinit__is_trivial = True
Methods
__init__
__init__(full_mbar: UnsafePointer[SharedMemBarrier, address_space=AddressSpace(3)], empty_mbar: UnsafePointer[SharedMemBarrier, address_space=AddressSpace(3)], warp_group_thread_idx: UInt) -> Self
Initialize ring buffer with barrier pointers.
__enter__
__enter__(mut self) -> Self
Context manager entry.
get_slot
get_slot[expected_bytes: Int = 0](mut self) -> UInt32
Producer waits for empty buffer slot and prepares for loading.
Returns:
enqueue_tile
enqueue_tile(mut self)
Producer signals for cp.async operations.
This handles the specific signaling pattern needed for cp.async:
- Signal async copy arrival
- Arrive at the barrier
- Advance to next stage
get_tile
release_slot
release_slot(mut self, read_idx: UInt32)
Consumer signals that buffer slot is empty.
arrive_empty_barriers
arrive_empty_barriers(self)
Helper to arrive at empty barriers during consumer initialization.
Was this page helpful?
Thank you! We'll create more content like this.
Thank you for helping us improve!