Asynchronous Copy, mbarrier & TMA

Recovered by disassembling ptxas v13.1.115 (CUDA 13.1) output with nvdisasm V13.1.115, cross-validated against the decoded per-architecture SASS instruction tables. Each PTX→SASS mapping below was produced by assembling a probe kernel for the relevant sm_ target and reading the emitted machine code. Facts and models only — the verbatim NVIDIA tables are not redistributed.

Three generations of GPU added a decoupled memory engine that moves data from global memory into shared memory without occupying the math datapath, and a shared-memory completion object (the mbarrier) that signals when the bytes have landed. This page recovers the full lowering of that subsystem from the PTX the programmer writes down to the SASS the GPU executes:

• Ampere cp.async — the LDGSTS instruction (global→shared), the commit/wait group mechanism (LDGDEPBAR + DEPBAR.LE), and the Ampere arrive (ARRIVES.LDGSTSBAR).
• The mbarrier object — its packed 64-bit shared-memory layout (phase, pending arrivals, expected arrivals, transaction-byte count) and the Hopper SYNCS family that operates on it.
• Hopper TMA (cp.async.bulk / cp.async.bulk.tensor) — the UBLKCP / UTMALDG / UTMASTG / UTMAREDG / UTMAPF engine, the 128-byte tensor-map descriptor, multicast, im2col, and the UTMACMDFLUSH drain.

The whole subsystem is compute-only: every instruction in it carries VALID_IN_SHADERS = (1<<ISHADER_TRAP) | (1<<ISHADER_CS) and is structurally illegal in any graphics shader (see Legality & Capabilities).

1. Ampere cp.async — LDGSTS

cp.async.<cache>.shared.global copies bytes from global memory straight into shared memory, never passing through a register. ptxas lowers it to a single LDGSTS (Load-Global-Store-Shared) instruction. Assembling the cache/size matrix for sm_80 and disassembling gives the exact form:

The mnemonic decoration order is fixed:

LDGSTS.E{.BYPASS}{.LTC64B|.LTC128B|.LTC256B}{.64|.128}{.ZFILL}

with the following semantics, each recovered from a one-variable probe:

• .E — extended (64-bit) global address. Always present; the global pointer occupies a register pair ([Rn.64]).
• Width (<empty>/.64/.128) — the per-thread byte count: 4, 8 or 16 bytes, taken directly from the PTX size operand. Sixteen-byte copies are the only width that accepts .cg.
• .BYPASS — the L1 cache-bypass form. cp.async.cg (cache-at-global-L2-only) maps to .BYPASS; cp.async.ca (cache-at-all-levels) leaves it off. The two are distinct cache-residency policies, not just a hint: .ca writes the line into L1 on its way to shared, .cg/.BYPASS does not.
• .ZFILL — out-of-bounds zero-fill. When the PTX carries the optional source-size operand N (the number of valid bytes when fewer than the copy width are in bounds), ptxas emits .ZFILL; the hardware copies N valid bytes and writes zeros for the remaining copy-size − N bytes of the shared destination. This is how cp.async boundary tiles are handled without a branch.
• .LTC64B/.LTC128B/.LTC256B — an L2-residency prefetch hint (L2::64B/128B/256B), independent of the L1 policy. It pulls the surrounding L2 sector group toward the SM so neighbouring copies hit warm.

A predicated cp.async lowers to a guarded @P LDGSTS …; the predicate gates the copy issue, and the .ZFILL semantics are independent of it.

LDGSTS carries the primary opcode byte 0xae. The four decoration groups map onto four independent encoding sub-fields, each pinned by an isolated probe:

The eviction class shares the same cop field (EF/EN/EL/LU/EU/NA) as an ordinary LDG (see Memory Model).

The async-copy group — commit and wait

cp.async copies are tracked as ordered groups on a single dedicated async scoreboard. There are two lowering rules, confirmed by assembling interleaved commits and waits:

LDGDEPBAR (Load-Global-Store-Shared Dep-Bar) takes no operands; it closes the current group and advances the group counter on async scoreboard SB0. DEPBAR.LE SB0, k blocks the warp until the number of still-outstanding copy groups on SB0 is ≤ k. So wait_group 0 / wait_all blocks until every committed copy has retired; wait_group 2 lets the two most-recent groups stay in flight. A three-group probe makes the counting explicit:

LDGSTS.E.128 …        ; group 0 body
LDGDEPBAR             ; commit group 0
LDGSTS.E.128 …        ; group 1 body
LDGDEPBAR             ; commit group 1
LDGSTS.E.128 …        ; group 2 body
LDGDEPBAR             ; commit group 2
DEPBAR.LE SB0, 0x1    ; wait_group 1 — tolerate ≤1 outstanding
DEPBAR.LE SB0, 0x0    ; wait_group 0 — drain all

All cp.async traffic funnels through SB0; the LDGDEPBAR/DEPBAR.LE pair is an explicit partial-wait that is byte-identical from SM80 through SM120 (see the dispatch-stall table in Execution Model). The DEPBAR encoding carries three fields: the scoreboard index, the count threshold k, and a .LE (less-than-or-equal) test flag; ptxas always sets .LE and the dedicated async-scoreboard index for the cp.async path.

Coupling a copy group to a barrier — ARRIVES.LDGSTSBAR

Instead of (or in addition to) the group counter, an Ampere cp.async batch can signal an mbarrier when it completes. cp.async.mbarrier.arrive[.noinc] .shared.b64 [bar] lowers to a dedicated arrive form:

LDGSTS.E.128 …                         ; the async copy
ARRIVES.LDGSTSBAR.64 [URZ+0x1000]      ; arm: signal [bar] when this copy lands

ARRIVES.LDGSTSBAR is the Ampere instruction that ties the LDGSTS completion to a shared-memory barrier: the warp arrives on the named mbarrier after the outstanding cp.async copies have written shared memory, not when the ARRIVES itself issues. This is the direct ancestor of Hopper's transaction-counting SYNCS.ARRIVE (§2). The .noinc variant arrives without incrementing the arrival count (it only registers the copy-completion dependency).

2. The mbarrier object

An mbarrier is a 64-bit object that lives in shared memory and tracks two things at once: a thread-arrival count (how many participants have arrived) and a transaction-byte count (how many bytes a pending async copy still owes). When both are satisfied the phase bit flips and any waiter on the old phase is released. The PTX type is .b64; the SASS that manipulates it differs by generation.

On Ampere the object is an opaque 64-bit shared word touched by ordinary shared-memory atomics — no dedicated barrier engine exists yet. The three operations lower to:

ATOMS.ARRIVE is a shared-memory atomic with an ARRIVE reduction op (the named-barrier atomic also used by bar-class arrivals); there is no transaction-byte channel, so an Ampere mbarrier only counts arrivals. Hopper replaces this with a purpose-built barrier engine, the SYNCS family, which adds the transaction-byte counter that TMA copies decrement (§3). The rest of this section describes that Hopper/Blackwell SYNCS form.

Packed 64-bit layout — recovered from mbarrier.init

mbarrier.init.shared.b64 [bar], count lowers to a single 64-bit shared store. Varying count and reading the init word out of the disassembly fixes the encoding exactly:

; mbarrier.init.shared.b64 [bar], C
UIADD3 t,  -C, 0x100000, URZ     ; t = 2^20 - C
USHF.L  hi, t, 0xb,  URZ         ; hi-word field = t << 11   (overall bit 43)
USHF.L  lo, t, 0x1,  URZ         ; lo-word field = t << 1    (overall bit 1)
SYNCS.EXCH.64 URZ, [bar], {lo,hi}   ; atomically write the packed init word

Both arrival fields are stored as the down-counting complement 2^20 − count, so a 20-bit field counts up toward 2^20 (zero crossing = satisfied) rather than down toward zero:

• The pending-arrival count occupies a 20-bit field starting at bit 1 of the word (t << 1).
• The expected-arrival count — the reset value reloaded after each phase flip — occupies a 20-bit field starting at bit 43 (t << 11 within the high 32-bit half).
• The remaining low bit (bit 0) is the phase bit; the high bits above the expected field hold the transaction-byte counter, which mbarrier.init leaves at zero.

mbarrier.inval.shared.b64 [bar] lowers to SYNCS.CCTL.IV [bar] — it invalidates the object's cache state so the shared slot can be reused for non-barrier data.

Arrive, expect-tx, and wait — the SYNCS family

Assembling each mbarrier op for sm_90a and sm_100a yields an identical mapping on both:

The TRANS64 decoration marks the instruction as operating on the 64-bit transaction-counting barrier object. The arrive modifiers encode which fields the instruction touches:

• SYNCS.ARRIVE.TRANS64 — a plain arrival; decrements the pending-arrival field by one (advancing the up-counter).
• .A1T0 (Arrive-1, Tx-0 set) — arrive.expect_tx: register one arrival and raise the transaction-byte count by N, so the barrier will not flip until N more bytes have been delivered by async copies. This is the producer half of a TMA pipeline (§3).
• .RED.A0TR (Reduce, Arrive-0, Tx-Raise) — the expect_tx-only form: add to the transaction-byte expectation without registering a thread arrival.
• SYNCS.PHASECHK.TRANS64{.TRYWAIT} — read the phase bit and compare it to the expected parity ph, returning a predicate. .TRYWAIT is the non-blocking try_wait (returns immediately with the current truth value, for spin loops); the plain form is test_wait. A waiter loops SYNCS.PHASECHK … @!P bra back until the phase flips.

Single-thread issue — ELECT

Operations that must run on exactly one thread of the warp (arming a barrier, issuing a bulk copy) are guarded by ELECT, the warp leader-election instruction:

@P0 ELECT P1, URZ, PT    ; P1 = (this lane is the elected leader)
@P1 …                    ; only the leader proceeds

ELECT returns a predicate true in exactly one active lane (and the elected lane id in a uniform register). ptxas wraps the descriptor build, the expect_tx arm, and the UBLKCP/UTMALDG issue in an ELECT guard so the single-issue semantics of the async engine are respected without the programmer writing the election by hand.

3. Hopper TMA — cp.async.bulk and cp.async.bulk.tensor

The Tensor Memory Accelerator (TMA) is a dedicated copy engine that moves bulk regions — contiguous byte ranges or multi-dimensional tensor tiles — between global and shared memory, addressed once per warp through a uniform-register path and completed through an mbarrier's transaction-byte counter. ptxas exposes two instruction sub-families: the UBLK* ops for untyped byte ranges and the UTMA* ops for tensor tiles described by a tensor-map descriptor.

Untyped bulk copies — UBLKCP / UBLKRED / UBLKPF

Assembling the cp.async.bulk matrix for sm_90a:

UBLKCP.<dst>.<src> names its direction by the destination then source memory space: .S.G is global→shared (the load), .G.S is shared→global (the store). The shared address fed to the engine is a cluster-mapped DSMEM address built from the CTA rank and the shared offset:

S2UR UR, SR_CgaCtaId          ; this CTA's rank within the cluster
ULEA UR_addr, UR, UR_off, 0x18 ; DSMEM addr = (rank << 0x18) | shared_offset

The << 0x18 (24-bit) shift packs the CTA rank into the high bits of the shared address, which is how a shared::cluster copy targets another CTA's shared memory inside the cluster. A shared→global UBLKCP.G.S or a UBLKRED is followed by UTMACMDFLUSH (op 0xb7), which drains the TMA command queue so the subsequent DEPBAR.LE SB0, 0x0 (cp.async.bulk.wait_group 0) observes completion. The .read wait variant (cp.async.bulk.wait_group.read) additionally emits CCTL.IVALL to invalidate the source so the buffer can be reused.

The bulk-copy opcode bytes (low byte of the 128-bit word) are UBLKCP = 0xba, UBLKPF = 0xbc; the issue is wrapped in an ELECT P0, URZ, PT (op 0x2f) leader guard so exactly one lane launches the engine.

Tensor copies — UTMALDG / UTMASTG / UTMAREDG / UTMAPF

cp.async.bulk.tensor.<rank>d copies a tile of a multi-dimensional tensor using a descriptor (the tensor-map) plus per-call tile coordinates. The matrix for sm_90a (identical on sm_100a):

The descriptor is read once into a uniform-register pointer (ULDC.64 UR, c[0x0][0x210] — the kernel-parameter bank on Hopper) and the instruction is issued under an ELECT guard so exactly one thread launches it:

@P0 ELECT P1, URZ, PT
    UTMALDG.2D [smem], [desc]        ; one elected thread issues the tile load

The decoration encodes the geometry directly:

• .2D / .3D / ….5D — tensor rank, from the PTX .<n>d modifier.
• .TILE (implicit) vs .IM2COL — the access pattern. IM2COL is the convolution-lowering mode that gathers an implicit im2col patch from a higher- rank tensor; it takes extra per-call offset operands (the convolution corner offsets) beyond the box coordinates.
• .MULTICAST — the cluster-broadcast load. A cp.async.bulk.tensor .multicast::cluster reads the tile once from global memory and writes it into the shared memory of every CTA selected by the 16-bit CTA-mask operand, saving global bandwidth across the cluster.
• .ADD / .MIN / … on UTMAREDG — the reduction operator applied when storing shared→global, the SASS half of cp.reduce.async.bulk.tensor.

Completion is signalled through the mbarrier whose address is passed in the PTX (mbarrier::complete_tx::bytes): the producer thread first arms the barrier with SYNCS.ARRIVE.TRANS64.A1T0 …, N (expecting N transaction bytes), then issues the UTMALDG; as the tile lands the engine decrements the barrier's transaction-byte counter, and when it reaches zero (and arrivals are satisfied) the phase flips, releasing the consumer's SYNCS.PHASECHK.TRANS64.TRYWAIT wait loop. A full producer→consumer pipeline disassembles to:

SYNCS.EXCH.64 URZ, [bar], packed        ; mbarrier.init
SYNCS.ARRIVE.TRANS64 RZ, [bar], 0x4000  ; arrive + expect_tx = 16384 bytes
@P0 ELECT P1, URZ, PT
    UTMALDG.2D [smem], [desc]            ; issue tile load (one thread)
WAIT:
SYNCS.PHASECHK.TRANS64.TRYWAIT P0, [bar], phase
@!P0 BRA WAIT                           ; spin until the tile has landed

The 128-byte tensor-map descriptor

The tensor-map is a 128-byte (1024-bit) opaque structure. ptxas does not build it on-device for static cases — it is normally produced host-side — but the tensormap.replace PTX ops do lower to plain shared loads/stores with bitfield masking, which exposes the byte offsets of each field. Assembling each tensormap.replace.tile.<field> in isolation for sm_90a and reading the STS offset recovers the layout:

So the descriptor encodes everything the engine needs to walk a tile: the global base and per-dimension global strides (in 16-byte units) and global extents (global_dim), the box dimensions of the tile to fetch, the element stride, the swizzle mode (the shared-memory bank-conflict-avoiding layout applied as the tile lands), the interleave layout, and the OOB-fill mode (whether out-of-bounds tile elements are zero-filled or take a constant — the tensor-tile analogue of cp.async's .ZFILL).

Publishing a descriptor that was modified on-device requires a proxy fence: tensormap.cp_fenceproxy … lowers to an ATOMG.E.EXCH.STRONG.GPU (a strong-scope exchange) that flushes the generic-proxy write of the descriptor so the TMA engine, which reads it through a different memory proxy, observes the new bytes. The acquire half — fence.proxy.tensormap::generic.acquire … — lowers instead to UTMACCTL.IV (op 0xb9), which invalidates the TMA's tensor-map proxy cache so a subsequent tile op re-reads the freshly published descriptor.

4. Named barriers and cluster barriers (for contrast)

The async barrier object (§2) is distinct from the classic named/CTA barrier and the cluster barrier, which share the BAR and UCGABAR opcode space:

BAR.SYNC is the whole-CTA convergence barrier; the .DEFER_BLOCKING modifier is the Hopper form that lets the warp defer the block so independent work can issue before the barrier actually stalls. BAR.ARV is the split-barrier arrive (no wait). UCGABAR_ARV/UCGABAR_WAIT are the cluster-wide (CGA) barrier arrive/wait, paired with CGAERRBAR (the cluster error-barrier that drains the cluster memory error pipe). These are count-based barriers with no transaction-byte channel — they synchronise threads, not async byte deliveries.

5. How ptxas sequences async ops against the scoreboard

Two orthogonal completion mechanisms cover the whole subsystem:

1. Group / scoreboard completion (cp.async, cp.async.bulk without an mbarrier). The copies are tracked as ordered groups on async scoreboard SB0; LDGDEPBAR (or the implicit bulk_group commit) advances the group counter and DEPBAR.LE SB0, k waits the count down. UTMACMDFLUSH first drains the TMA command queue so the DEPBAR sees the right count. This is the low-overhead path when the consumer is in the same warp and a simple FIFO ordering is enough.

2. Transaction-byte completion (cp.async.bulk[.tensor].mbarrier… and the Ampere cp.async.mbarrier.arrive). The copy carries the address of an mbarrier; the producer arms the expected byte count (SYNCS.ARRIVE.TRANS64.A1T0, or on Ampere the ARRIVES.LDGSTSBAR form), the engine decrements the byte counter as data lands, and a separate consumer warp is released by the phase flip it observes through SYNCS.PHASECHK. This decouples the producer and consumer entirely — the consumer never touches the scoreboard, only the barrier object.

ptxas chooses the mechanism from the PTX form (a .mbarrier::complete_tx::bytes qualifier selects path 2, a .bulk_group/.commit_group selects path 1), wraps single-issue ops in ELECT, inserts UTMACMDFLUSH before any group-wait that follows a shared→global bulk op, and emits the proxy fence (ATOMG.E.EXCH.STRONG.GPU) for any on-device descriptor it must publish. The async scoreboard SB0 and the per-instruction wait-mask in the scheduling-control word are the two hardware resources it allocates; everything else — the barrier object, the transaction counter, the DSMEM cluster mapping — is built in shared memory by the instruction stream itself.

Architecture coverage

On Blackwell (sm_100a) the mbarrier and TMA instruction families are encoded identically to Hopper; the visible delta in the disassembly is that ordinary LDGSTS switches to the descriptor-addressed form LDGSTS.E … [Rd], desc[UR][Rsrc.64], where the global access goes through a uniform-register memory descriptor rather than a bare 64-bit pointer — the same descriptor-addressing change that the rest of the Blackwell global-memory path adopts.

Cross-References

• Execution Model — the async scoreboard SB0, the DEPBAR.LE wait-down, and the per-warp dispatch rule that sequences these ops.
• Memory Hierarchy & Ordering — how the cp.async/TMA path relates to the MEMBAR scope ladder and the LSU/coalescer model.
• Legality & Capabilities — the VALID_IN_SHADERS = TRAP|CS gate that makes the whole subsystem compute-only.
• Architecture Evolution — the SM75→SM121 ISA diff that places LDGSTS at Ampere and the TMA/SYNCS/cluster families at Hopper.
• GMMA/WGMMA Pipeline — the tensor-core consumer that these async copies feed.