Synchronization

Barriers

For the concepts behind these calls — asynchronous barrier semantics, the mbarrier phase model, named vs memory barriers, and producer-consumer patterns — see Barrier Support in TLX.

Allocates buffer in shared memory and initialize mbarriers with arrive_counts.

Input:

Use scope="cluster" for barriers that require cross-CTA initialization visibility. With multiple physical CTAs, if any eligible kernel-entry barrier has cluster scope, the compiler batches all eligible entry barrier initializations, including CTA-scoped barriers, before TMEM allocation. Immediately after the last initialization it inserts one initialization fence followed by a relaxed cluster arrive and cluster wait, independently of which operations consume the barriers. This works with either logical multi-CTA distribution or a physical multi-CTA cluster. A single physical CTA needs no cluster rendezvous; scope never changes the launch dimensions or barrier layout.

Allocate cluster-scoped barriers unconditionally in the kernel prologue, before their use. Consumers may be inside loops or warp-specialized tasks, but initializing these barriers inside branches, loops, or tasks is not supported. The default preserves existing behavior, including legacy compiler synchronization for recognized paired-MMA and other consumers.

    # Requires a two-CTA launch; scope does not create the cluster.
    bars = tlx.alloc_barriers(num_barriers=1, scope="cluster")
    # Initialization is ready for cross-CTA use; no manual init fence is needed.
    if tlx.cluster_cta_rank() == 0:
        tlx.barrier_arrive(bar=bars[0], remote_cta_rank=1)

Initialization scope does not change arrive_count, register expected transaction bytes, or wait for asynchronous completion. Keep data-ready handshakes, proxy fences, and DSM lifetime synchronization required by the algorithm. The automatic initialization rendezvous does not publish later generic stores or synchronize later pipeline phases.

Wait until the mbarrier phase completes

Perform the arrive operation on an mbarrier. If remote_cta_rank is provided, signals the barrier in the specified remote CTA's shared memory (useful for multi-CTA synchronization). On AMD this requires arrive_count == 1.

Wait until num_threads total threads have reached the specified named mbarrier phase.

Signal arrival at a named mbarrier.

For both APIs, num_threads is the total number of threads required to flip the barrier phase: num_waiting_threads + num_arriving_threads. Wait and Arrive calls for the same barrier phase must use the same value.

The backend lowers these to PTX bar.sync/bar.arrive (the .aligned forms) whenever the barrier provably executes warp-uniformly -- i.e. all threads in a warp execute the same barrier instruction -- and to the non-aligned barrier.sync/barrier.arrive otherwise. Warp-level divergence (warp specialization, branches on warp IDs) is compatible with .aligned; divergence on thread IDs or on data that can vary by thread within a warp is not.

Signal a barrier of an expected number of bytes to be copied.

Cluster initialization ownership

scope="cluster" declares an initialization requirement, not a new barrier storage type or a user-facing region API. The compiler constructs a designated kernel-entry prologue segment. With multiple physical CTAs, one cluster-scoped entry barrier triggers batching of all eligible entry initializations, including CTA/default-scope barriers. A single physical CTA needs no cluster rendezvous.

The initializing thread publishes mbarrier.init with an initialization release fence. The relaxed cluster arrive/wait prevents remote use before the owner is ready. The entire batch precedes TMEM setup, and hoisting precedes allocation analysis so shared-memory liveness sees the actual schedule. Initialization must be unconditional at entry; branch, loop, task, reinitialization, and later-phase cluster initialization are unsupported. Consumers may still be inside loops or WS tasks.

The compiler reuses existing Triton NVIDIA fence/arrive/wait IR and lowering; it does not call the Gluon Python API. TLX still defines the entry bootstrap ownership and ordering. Its exact coveredInitOps set tells the later cross-CTA initialization analysis to skip only that already-protected batch; other initialization and runtime dependency analysis remain active.

For WS bootstrap, worker warps arrive once in the header while default warps execute initialization, fence, arrive, and wait. Only the designated bootstrap operations bypass ordinary synthetic-WS wrapping, avoiding duplicate worker arrivals. Later synchronization keeps its normal lowering. This participation mechanism does not infer application data-ready or pipeline-phase requirements.

The default CTA scope retains legacy consumer-based synchronization. Do not remove required application handshakes merely because initialization now has cluster scope. For more background, see Barrier Support in TLX.

Scheduling barriers

Prevents selected AMD machine-instruction classes from crossing a source boundary. It is a scheduling marker, not a workgroup barrier or memory fence.

Memory fences

Examples:

  tlx.local_store(smem_buf, data)
  tlx.fence("async_shared")
  tlx.async_descriptor_store(desc, smem_buf, offsets)

  tlx.fence("async_shared_cluster")

Manual initialization example:

  bars = tlx.alloc_barriers(num_barriers=1, arrive_count=1)
  tlx.fence_mbarrier_init_cluster()
  tlx.cluster_barrier()

  # now bars is ready for cross CTA use
  tlx.barrier_arrive(bar=bars[0], remote_cta_rank=1)

NVIDIA DSM Memory Proxies

An operation's asynchronous completion does not determine its memory proxy. For NVIDIA SM90+ cluster shared memory, the following APIs use distinct paths:

st.async is a non-blocking store in the generic proxy, including its mbarrier completion access. It is not a shared-memory bulk copy. In contrast, cp.async.bulk reads the source buffer through the async proxy. See the NVIDIA PTX ISA for the instruction semantics.

Before async_remote_shmem_copy reads a source buffer written through the generic proxy, establish the handoff with tlx.fence("async_shared_cluster"). If another CTA produced the data, first establish completion and cross-CTA visibility, then issue the proxy fence in the CTA performing the bulk copy.

These mechanisms have different responsibilities:

mbarrier. Waiting for st.async to finish does not itself establish generic-to-async proxy ordering.

synchronization. It does not replace the proxy fence.

and async proxies for cluster shared memory. It does not wait for the bulk copy to complete; keep the source and destination alive until completion.

initialization synchronization. Its initialization fence and bootstrap rendezvous do not replace a proxy fence for later payload writes.