The CUDA Handbook ▸ Chapter 5. Memory

5.9 Intra-Kernel Asynchronous Memcpy

The shared-memory kernels of Section 5.6 load their data synchronously: every thread reads from global memory, writes the value into shared memory, and the block waits at __syncthreads() before any computation begins. That load phase carries two costs. The data takes a detour through registers – from global memory to a register to shared memory – and, because the loads and the barrier sit on the critical path, the SM stalls on global-memory latency with no other work scheduled to cover it.

Older hardware had essentially one way to absorb that latency: thread-level parallelism. Keep enough warps resident and the scheduler can always switch to one that is ready to run while others wait on memory, so the stalls overlap rather than accumulate. But the depth of that cover is the occupancy (Section 7.3), and occupancy is exactly the budget a reuse-heavy kernel spends elsewhere: the register blocking and large shared-memory tiles that raise arithmetic intensity draw on the same registers and shared memory that limit the number of warps that can stay resident. The harder a kernel works to reuse data, the lower its occupancy, and the fewer warps it has left to hide its own loads – so covering latency and maximizing reuse are competing objectives.

The Ampere architecture (SM 8.0) added an asynchronous copy instruction, cp.async, that addresses both costs. It moves bytes directly from global memory into shared memory, bypassing the register file, and it does not block the warp that issues it: the copy proceeds in the background while the warp does other work, and the kernel waits for completion only when it needs the data. The overlap now comes from the copy pipeline rather than from a surplus of warps, which is what lets a register-hungry kernel hide its loads at the low occupancy that reuse demands. Asynchronous copies are committed in groups called stages, and a kernel waits on stages in the order it committed them.

CUDA exposes cp.async at two levels. The low-level intrinsics in <cuda_pipeline.h> – __pipeline_memcpy_async(), __pipeline_commit(), and __pipeline_wait_prior() – issue 4-, 8-, or 16-byte copies and track the stages by hand. The libcu++ standard library wraps them in a portable, composable form: cuda::memcpy_async() copies a span of any size, coordinated by a cuda::pipeline or a cuda::barrier. The underlying pattern is the same either way – stage the copy, do unrelated work, then wait:

__shared__ float tile[TILE];

// Stage a copy of one element per thread from global into shared memory.
__pipeline_memcpy_async( &tile[threadIdx.x], &gmem[base + threadIdx.x], sizeof(float) );
__pipeline_commit();

// ... independent work that does not touch tile[] can run here ...

__pipeline_wait_prior( 0 );   // the staged copy has completed
__syncthreads();              // tile[] is now visible to the whole block

Staging a copy only to wait on it immediately merely relocates the latency. The technique helps only when a kernel double-buffers: while the arithmetic units consume the slab in one shared-memory buffer, the next slab streams asynchronously into a second buffer, so copy and computation proceed together instead of the copy stalling the pipeline. The interlock is essential – a consumer must never read a buffer before its copy has landed – and __pipeline_wait_prior() is what enforces the ordering.

Overlap is not free throughput, and asynchronous copy does not always help. It improves performance only when two conditions hold at once: the kernel has enough compute to hide the copy behind, and the copy was exposed on the critical path to begin with. A kernel whose loads already overlap its arithmetic gains nothing by moving them off the issuing warp.

Chapter 17 (Matrix Multiplication) measures both cases. Its Tensor Core GEMM (Section 17.6, Listing 17-6) alternates two shared-memory slabs exactly as described here: each step waits on the slab it is about to read, issues the copies for the next slab, then feeds the current one to the Tensor Cores. There the overlap is worth measurable throughput, because the Tensor Cores retire a slab fast enough to leave the global load exposed. The same cp.async double-buffering added to that chapter’s FP32 register-blocked kernel makes no difference – that kernel already hides its loads behind its multiply-adds, so no exposed latency remains to overlap.

On Hopper, the Tensor Memory Accelerator (TMA) extends asynchronous copy to bulk, multidimensional transfers. A single instruction moves a whole tile between global and shared memory from a descriptor prepared ahead of time, sparing the threads the address arithmetic that cp.async still costs. TMA works with the thread block clusters and distributed shared memory of Section 7.8, and it drives the fastest Hopper matrix multiplications.