Mojo module
This module provides GPU synchronization primitives and barriers.
The module includes:
- Block-level synchronization barriers (barrier())
- Warp-level synchronization (syncwarp())
- Memory barriers (mbarrier) for NVIDIA GPUs
- Instruction scheduling controls for AMD GPUs
- Asynchronous copy and bulk transfer synchronization
The synchronization primitives help coordinate execution between threads within thread blocks and warps, and manage memory consistency across different memory spaces.
comptime values
MaxHardwareBarriers
comptime MaxHardwareBarriers = 16
Maximum number of hardware barriers available per block.
Structs
-
AMDScheduleBarrierMask: Represents different instruction scheduling masks for AMDGPU scheduling instructions.
Functions
-
async_copy_arrive: Makes a memory barrier track all prior async copy operations from this thread. -
barrier: Performs a synchronization barrier at the block level. -
cp_async_bulk_commit_group: Commits all prior initiated but uncommitted cp.async.bulk instructions into a cp.async.bulk-group. -
cp_async_bulk_wait_group: Waits for completion of asynchronous bulk memory transfer groups. -
mbarrier_arrive: Signal thread arrival at a shared memory barrier. -
mbarrier_arrive_expect_tx_relaxed: Configure a shared memory barrier to expect additional async transactions. -
mbarrier_arrive_expect_tx_shared: Configure a shared memory barrier to expect additional async transactions. -
mbarrier_init: Initialize a shared memory barrier for synchronizing multiple threads. -
mbarrier_test_wait: Test if all threads have arrived at the memory barrier. -
mbarrier_try_wait_parity_shared: Wait for completion of a barrier phase with timeout. -
named_barrier: Performs a named synchronization barrier at the block level. -
named_barrier_arrive: Arrives at a named synchronization barrier at the block level. -
s_waitcnt: Performs ans_waitcntinstruction with the specified counters on AMD GPUs. -
s_waitcnt_barrier: Performs ans_waitcntfollowed by a barrier on AMD GPUs. -
schedule_barrier: Controls instruction scheduling across a barrier point in AMD GPU code. -
schedule_group_barrier: Controls instruction scheduling across a barrier point in AMD GPU code by creating schedule groups. -
syncwarp: Synchronizes threads within a warp using a barrier. -
umma_arrive_leader_cta: Signal arrival at the barrier to the leader CTA of the pair.