copy_async

Asynchronous copy. Every variant emits the transfer’s issue instruction; the caller supplies the matching completion protocol. ldgsts uses cp.async commit/wait, TMA loads and distributed-shared-memory copies signal an mbarrier, TMA stores use the bulk async-group commit/wait operations, and tensor-memory paths use their tcgen05 commit/wait operations. Selection is by the source/destination memory pair and scope. CUDA currently registers six variant names; the TMA page below covers two of them.

Variant	Pair	Prio	Issue instruction
copy_async → ldgsts	global → shared	20	cp.async (LDGSTS), per-thread vectorized
copy_async → tma_auto / tma_explicit (tma_auto / tma_explicit)	global ↔ shared	10	cp.async.bulk.tensor (TMA, descriptor-driven, single-thread)
copy_async → dsmem	shared → shared (cross-CTA)	10	cp.async.bulk shared::cluster (mapa remote address)
copy_async → smem->tmem (tcgen05.cp) (smem->tmem)	shared → tmem	10	one of the registered tcgen05.cp shapes (matrix-descriptor driven)
copy_async → tmem<->local (tcgen05.ld/st) (tmem<->local)	tmem ↔ register	10	tcgen05.ld / tcgen05.st (warpgroup, atom-matched)

• copy_async → ldgsts
• copy_async → tma_auto / tma_explicit
• copy_async → dsmem
• copy_async → smem->tmem (tcgen05.cp)
• copy_async → tmem<->local (tcgen05.ld/st)