Mojo struct

RingBuffer

struct RingBuffer[num_pipeline_stages: Int, num_consumers: Int, cluster_size: Int]

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)]):

Implemented traits

AnyType, UnknownDestructibility

Aliases

`delis_trivial`

alias __del__is_trivial = UnsafePointer[SharedMemBarrier, address_space=AddressSpace(3)].__del__is_trivial

Methods

`init`

__init__(out self, full_mbar: UnsafePointer[SharedMemBarrier, address_space=AddressSpace(3)], empty_mbar: UnsafePointer[SharedMemBarrier, address_space=AddressSpace(3)])

Initialize ring buffer with barrier pointers.

`init_barriers`

init_barriers(self, thread_idx: UInt)

Initialize the ring buffer barriers.

`init_barriers_cpasync`

init_barriers_cpasync(self, thread_idx: UInt)

Initialize barriers for cp.async variant.

`producer_wait_for_empty`

producer_wait_for_empty[expected_bytes: Int](self, write_idx: UInt32, pipeline_state: PipelineState[num_pipeline_stages])

Producer waits for empty buffer slot and prepares for loading.

`producer_signal_full`

producer_signal_full(self, write_idx: UInt32, mut pipeline_state: PipelineState[num_pipeline_stages])

Producer signals that buffer is full and advances to next stage.

`consumer_wait_for_full`

consumer_wait_for_full(self, read_idx: UInt32, pipeline_state: PipelineState[num_pipeline_stages])

Consumer waits for full buffer slot.

`consumer_signal_empty`

consumer_signal_empty(self, read_idx: UInt32, warp_group_thread_idx: UInt)

Consumer signals that buffer slot is empty.

`consumer_advance`

consumer_advance(self, mut pipeline_state: PipelineState[num_pipeline_stages])

Consumer advances to next pipeline stage.

`arrive_empty_barriers`

arrive_empty_barriers(self, warp_group_thread_idx: UInt)

Helper to arrive at empty barriers during consumer initialization.

Fields​

Implemented traits​

Aliases​

__del__is_trivial​

Methods​

__init__​

init_barriers​

init_barriers_cpasync​

producer_wait_for_empty​

producer_signal_full​

consumer_wait_for_full​

consumer_signal_empty​

consumer_advance​

arrive_empty_barriers​