Why architecture matters here
The architecture matters because most heavy GPU kernels are memory-latency-bound long before they are compute-bound, and latency is not fixed by bandwidth. You can have terabytes per second of HBM bandwidth and still stall, because a single dependent load still takes hundreds of cycles to return, and if the very next instruction needs that data, the warp waits. Tiling into shared memory reduces how many times you pay that latency, but it does not by itself hide the latency you still pay on each tile load. Asynchronous copy is the mechanism that converts 'wait for the tile, then compute' into 'compute this tile while the next one loads,' which is the difference between running at a fraction of peak and running at the roofline.
It matters because the naive synchronous copy actively sabotages occupancy, the other latency-hiding tool. A global -> shared load on older paths is really two instructions with a register in between, and that register is live for the entire hundreds-of-cycles load. A kernel loading a large tile ties up a whole block of registers purely as a data conduit, which caps how many warps fit per multiprocessor, which reduces the scheduler's ability to hide latency with other warps. So the old path makes you choose between register-hungry tiles and high occupancy. Register-free asynchronous copy removes the conduit entirely: bytes flow straight to shared memory, registers stay available for accumulators, and occupancy is no longer collateral damage of the load path.
The architecture also matters because it makes latency hiding explicit and controllable rather than emergent. With pure occupancy-based hiding, whether latency is actually hidden depends on the scheduler happening to have a ready warp — it works until register or shared-memory limits cut occupancy, and then it silently stops working. A software pipeline with a fixed number of buffer stages hides a known, designed amount of latency regardless of occupancy: if you prefetch three stages ahead, you tolerate three tiles' worth of load latency by construction. The performance becomes predictable and tunable — the stage count is a dial you turn against the shared-memory budget — instead of an accident of how the compiler allocated registers.
Finally, it matters because it is the foundation every high-performance library kernel is built on. The fast GEMM, attention, and convolution kernels in cuBLAS, CUTLASS, and FlashAttention are, at their core, multi-stage asynchronous-copy pipelines: prologue to fill the stages, a main loop that overlaps Tensor-Core math with the next tile's async load, epilogue to drain. Understanding this pattern is understanding why those kernels reach near-peak throughput and why a hand-written kernel that copies synchronously, or pipelines too shallowly, leaves most of the hardware on the table. It is the single most important structural idea in memory-bound GPU code.
The architecture: every piece explained
Top row: the register-free load into a staged buffer. The input lives in global memory — the large, high-bandwidth, high-latency HBM that holds the operands. Instead of a synchronous load that stages through a register, the kernel issues a cp.async: an instruction that copies a chunk (ideally a 128-bit aligned vector per thread) directly from global memory into shared memory, bypassing the register file, and returns immediately without waiting for the data. The destination shared memory is organized as a multi-stage buffer — several tile-sized slots — so that while one slot is being consumed by compute on tile k (Tensor Cores for matmul, or ALUs for elementwise work), other slots can be filling from in-flight asynchronous copies. The copy and the compute touch different buffer stages, so they proceed concurrently.
Middle row: the pipeline structure that turns concurrency into overlap. A copy alone does not hide latency; you have to arrange the loop so the copy is issued far enough ahead of its use. The prologue primes the pipeline by issuing asynchronous copies for the first N-1 stages before the main loop begins — filling the buffer ahead of any compute. Each batch of copies is grouped with a commit, and cp.async.wait_group lets the kernel block only on the oldest outstanding group — the tile it is about to consume — while newer copies keep flying. The steady state is the heart: in each iteration the kernel computes on the tile whose copy has completed and, in the same iteration, issues the asynchronous copy for a tile several steps ahead, so load and compute overlap continuously. The epilogue drains the final stages, computing on the tiles still in the buffer after the last copy has been issued.
Bottom-left: the payoff on latency. Because the copy for tile k+1 (and k+2, up to the stage depth) is issued while the kernel computes on tile k, the hundreds of cycles that load takes are spent doing useful arithmetic rather than waiting. The latency is hidden behind compute, so the multiprocessor stays busy and throughput approaches the limit set by whichever of memory bandwidth or compute is the true bottleneck — not by memory latency.
Bottom-right and ops: the freed budget and the levers. Because cp.async never stages through registers, the occupancy is freed from the register pressure a synchronous copy would impose — those registers go to accumulators, raising how many warps and how large a tile fit. The ops strip names the tuning surface: the stage count traded against the finite shared-memory budget (more stages hide more latency but consume more shared memory, which can itself cut occupancy), copies kept 128-bit aligned so each async transaction moves a full vector efficiently, the shared-memory layout arranged to be bank-conflict-free so the compute reads are not serialized, and the wait-group depth set so the kernel waits on exactly the right tile and no earlier.
End-to-end flow
Trace a tiled matrix-multiply kernel that computes a block of the output by streaming K-dimension tiles of its two inputs through a three-stage shared-memory pipeline, from the first prefetch to the final accumulate.
Prologue — fill the pipeline: before any math, the kernel issues asynchronous copies for the first two K-tiles of both operands into shared-memory stages 0 and 1, committing each as its own group. It does not wait for them yet; the copies are in flight over HBM while the kernel finishes setting up accumulators in registers. The third stage (stage 2) is left empty, ready to receive the next prefetch. At this point two tiles' worth of load latency is already being absorbed by the setup work and by the copies overlapping each other.
Steady state — overlap load and compute: the main loop over K-tiles begins. At the top of iteration k the kernel calls cp.async.wait_group to block until the oldest outstanding copy group — the tile it is about to multiply — has landed in shared memory, then a barrier makes it visible to the whole block. It immediately issues the asynchronous copy for tile k+2 into the stage that tile k-1 just vacated, committing that group so it flies while the math runs. Only then does it feed tile k's operands from shared memory into the Tensor Cores, accumulating into registers. Because the copy for k+2 was launched before the compute on k, and the copy for k+1 is already in flight, the multiprocessor never stalls waiting for memory — every tile's load latency is hidden behind the arithmetic of the tiles ahead of it.
Bank-conflict-free reads: as the Tensor Cores consume tile k, they read the shared-memory buffer in a pattern the kernel deliberately laid out (often with padding or a swizzle) so that the 32 lanes of a warp hit 32 distinct shared-memory banks. Without that layout, multiple lanes would collide on the same bank and the reads would serialize, throttling the compute the pipeline worked so hard to feed. The async copy filled the buffer efficiently; the conflict-free layout ensures the compute drains it just as efficiently.
Wait-group discipline: the kernel keeps exactly N-1 copy groups outstanding at all times by waiting for the group that is N-1 behind the one just issued. Waiting too aggressively (draining all groups every iteration) would collapse the overlap back into a synchronous load; waiting too loosely (never gating) would let compute race ahead of data that has not arrived, reading garbage. The wait-group depth is what holds the pipeline at exactly the right number of in-flight stages.
Epilogue — drain the tail: after the loop issues its last prefetch, two tiles remain in flight and unconsumed. The epilogue waits for and computes on those final stages one by one, with no new copies to issue, flushing the last of the accumulation. The kernel then writes the completed output block from registers back to global memory. Across the whole run, the HBM was read at close to bandwidth and the Tensor Cores were rarely idle — the memory latency that would have dominated a synchronous version was spent, tile after tile, doing math.