IMPORTANT: To view this page as Markdown, append `.md` to the URL (e.g. /max/get-started.md). For the complete documentation index, see llms.txt.
Skip to main content
For the complete documentation index, see llms.txt. Markdown versions of all pages are available by appending .md to any URL (e.g. /max/get-started.md).

Mojo module

pipeline_backend

Hardware synchronization backends for ProducerConsumerPipeline.

ProducerConsumerPipeline coordinates warp-specialized producer and consumer warps through a ring of shared-memory slots. Each slot has two signals:

  • full: the producer raises it when a slot holds fresh data.
  • empty: the consumer raises it when a slot has been drained and may be refilled.

The protocol (which slot, which lap around the ring, who waits on whom) is hardware-neutral and lives in ProducerConsumerPipeline. The mechanism used to raise and wait on a signal is hardware-specific and lives behind the PipelineBackend trait defined here:

  • NvidiaMbarBackend uses NVIDIA mbarrier objects (SharedMemBarrier).
  • A future non-NVIDIA backend (see KERN-2625) will use shared-memory atomic counters, which have no phase/parity concept.

Phase / lap convention​

The pipeline hands each backend a monotonically increasing phase (the lap number for a given slot, incremented each time the ring wraps). Backends map it onto their substrate:

  • NvidiaMbarBackend reduces it to the single parity bit phase & 1 that mbarrier.try_wait.parity expects. The parity sequence is identical to the historical ^= 1 toggle, so this backend is bit-for-bit equivalent to the previous hardcoded implementation.
  • A counter-based backend can derive the absolute threshold a slot must reach from phase directly, so it needs no per-slot mutable wait state.

comptime values​

MbarPtr​

comptime MbarPtr = Pointer[SharedMemBarrier, MutUntrackedOrigin, address_space=AddressSpace.SHARED, _safe=False]

Structs​

Traits​

Edit this page

Was this page helpful?

Models

Cookbook

Puzzles

Blog

Community

Contact