Asynchronous Global to Shared Copy Memory

Copy from global memory into shared or local memory without staging the data through registers.

Vendor-Neutral Definition

A DMA-style transfer issued by the thread and completed later, so the wait can be deferred and overlapped with unrelated arithmetic.
⚠️ Cross-vendor mappings describe semantic relationships and are not guaranteed one-to-one compiler translations.
equivalent with restrictions Both move data from global memory to on-chip scratchpad without occupying vector registers, and both are asynchronous, but they differ in how completion is tracked.
Property NVIDIA PTX AMDGPU
Instruction cp.async global_load_lds_dword
ISA Layer Virtual Native
Data Types - -
Version / Target Introduced PTX ISA 7.0 -

Important Differences

  • PTX cp.async is grouped into commit batches and waited on with cp.async.wait_group, so a kernel can have several batches in flight and wait for the oldest. AMDGPU tracks the load with the ordinary vmcnt counter and waits with s_waitcnt.
  • cp.async arrived with Ampere (sm_80); the AMDGPU LDS-direct loads are older and present across CDNA parts, so the portability constraint runs in opposite directions on the two vendors.
  • Both bypass the vector register file, which is the point: on a software-pipelined GEMM this is what frees registers for accumulators.

Source Evidence

Verification method: documentation