Inline PTX: cp.async Double Buffering

Lesson

The harness fixes the shape: 256 threads per block, four tiles of 1,024 floats per block, and a grid derived from N. asyncTransform is the arithmetic. The starter loads each tile through registers, waits at a barrier, computes, waits again, and only then starts the next tile, so the copy hardware and the SM take turns. The memory pipelines chapter overlapped them with cuda::memcpy_async and a thread-scoped pipeline. This lab writes the instructions that API lowers to.

cp.async.cg.shared.global [dst], [src], 16; copies 16 bytes from global to shared memory without passing through a register, and the thread keeps running while the bytes travel. The .cg qualifier caches at L2 only and accepts 16-byte copies; .ca accepts 4, 8 or 16 and keeps the line in L1 as well. dst is a 32-bit shared-space address: take __cvta_generic_to_shared of the shared pointer, cast it to unsigned, and bind it with "r". src is a global-space address from __cvta_generic_to_global, bound with "l". Both sides must be 16-byte aligned; the tiles are declared alignas(128) and each chunk is four floats, so every chunk address is aligned.

Copies complete in groups. cp.async.commit_group; closes every copy this thread has issued since the last commit into one group. cp.async.wait_group N; blocks until at most N of this thread's groups are still in flight, so wait_group 0 means all of them have landed. The wait covers only the calling thread's own copies; a block barrier after it is what makes every thread's chunks visible to the block.

The pipeline is then a few lines. Prime: every thread issues its share of tile 0 into buffer 0 and commits. Each stage: if a next tile exists, issue it into the other buffer, commit, and wait_group 1, which returns once the older group, the current tile, has landed while the newer one may still be flying; on the last stage, wait_group 0. Barrier, compute from the current buffer, barrier again. That second barrier matters: the prefetch two stages later refills this buffer, and no thread may still be reading it. The chunk loop steps by blockDim.x, so with 256 threads and 256 chunks per tile each thread issues one copy per tile.

In the PTX pane, find your cp.async statements between the inline asm markers; the listing captured from the run holds the same instructions without the comments. In the SASS, LDGSTS is the global-to-shared copy, and the group waits appear as LDGDEPBAR and DEPBAR.LE. If you ran the memory pipelines lab, compare the two listings: cuda::memcpy_async produced the same instructions. The harness times your kernel against the synchronous baseline in the same binary and prints the speedup; correctness is graded for every size, and no uncalibrated threshold is required to pass.

What you will learn

• Issue a 16-byte asynchronous global-to-shared copy with cp.async and bind its shared and global addresses correctly.
• Close copies into groups and wait for the oldest group while a newer one is still in flight.
• Place the two barriers that make a double-buffered tile loop correct: one after the wait, one after the compute.
• Recognize cp.async, LDGSTS and the dependency barriers in a captured listing.

Starter code

    __shared__ alignas(128) float tile[kAsyncTileElements];
    int blockBase = static_cast<int>(blockIdx.x) * kAsyncElementsPerBlock;

    for (int stage = 0; stage < kAsyncTilesPerBlock; ++stage) {
        int tileBase = blockBase + stage * kAsyncTileElements;
        for (int i = static_cast<int>(threadIdx.x); i < kAsyncTileElements; i += blockDim.x) {
            tile[i] = input[tileBase + i];
        }
        __syncthreads();
        for (int i = static_cast<int>(threadIdx.x); i < kAsyncTileElements; i += blockDim.x) {
            output[tileBase + i] = asyncTransform(tile[i]);
        }
        __syncthreads();
    }