Asynchronous Prefetch: Double-Buffer Shared Tiles

Experimental — correctness is graded; the performance target awaits GPU calibration. The starter waits for each global-to-shared tile load before doing any arithmetic. Build a two-stage cuda::pipeline so tile n+1 moves into the unused shared buffer while the block computes tile n.

Languages: CUDA C++, Rust

Lesson

On compute capability 8.0 and newer, aligned global-to-shared cuda::memcpy_async operations can use hardware asynchronous copies that avoid an intermediate register. The performance benefit comes from overlap: issuing an async copy and immediately waiting for it is still serial execution.

SM global-to-shared asynchronous copies differ from host-device DMA engines. Synchronize producer/consumer stages and validate tails before interpreting overlap.

What you will learn

Create a thread-scoped pipeline: each thread tracks only its own copies, so no shared pipeline state is needed.
Prime the pipeline before entering the consume loop.
Prefetch the next tile into the opposite shared buffer while computing the current tile.
Pair producer and consumer operations without reusing a live stage.
Keep every block thread converged around pipeline waits and block synchronization.

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();
    }