CUDA Pipelines

Global memory accesses can take hundreds of cycles, making latency hiding one of the primary challenges in kernel design. An MMA tile may finish in less than ten cycles while a global load still takes hundreds. For tensor-core kernels, staging is usually the bottleneck. Cycle counts aren't officially documented; see this forums thread for back-of-envelope WGMMA timing (~4 vs ~18 cycles for MACs vs shared-memory reads).

CUDA pipelines, built on Ampere's asynchronous copy instructions, overlap data movement with computation. The API is simple; the hard part is choosing enough pipeline depth to cover memory latency without sacrificing occupancy through extra shared memory.

Motivation

A common pattern in CUDA kernels is tiled processing through shared memory:

for (int tile = 0; tile < num_tiles; tile++) {
    load_tile_from_global_to_shared();

    __syncthreads();

    compute_on_tile();

    __syncthreads();
}

While this approach improves memory locality, it also creates a strict dependency chain. Computation cannot begin until the entire tile has been loaded into shared memory. As a result, global memory latency becomes visible to the execution pipeline.

Historically, loading data into shared memory required a two-step process:

Global Memory
      ↓
   Register
      ↓
 Shared Memory

Every thread first loads data into a register and then stores it into shared memory. Besides the latency cost, this consumes registers that could otherwise be used for computation.

Ampere introduced hardware support for asynchronous copies from global memory directly into shared memory through the cp.async instruction. CUDA exposes this functionality through cuda::memcpy_async and cuda::pipeline, enabling software pipelining of memory transfers and computation.

The goal is simple: while one tile is being processed, the next tile should already be moving through the memory hierarchy.

The Pipeline Model

CUDA pipelines follow a producer-consumer model.

Instead of processing tiles sequentially, data movement and computation are overlapped:

         time ─────────────────────────→

Copy     │ Tile 0 │ Tile 1 │ Tile 2 │ Tile 3 │
Compute  │        │ Tile 0 │ Tile 1 │ Tile 2 │

The pipeline API exposes several key operations:

pipeline.producer_acquire();

cuda::memcpy_async(...);

pipeline.producer_commit();

pipeline.consumer_wait();

compute();

pipeline.consumer_release();

Conceptually, a pipeline behaves as a ring buffer of stages stored in shared memory. The producer side fills future stages with incoming data while the consumer side processes previously loaded stages.

Under the Hood: cp.async

Ampere implements the direct global-to-shared path in hardware via PTX instructions such as:

cp.async.ca.shared.global

Several benefits emerge:

• Reduced register pressure: Since data no longer passes through temporary registers, kernels often require fewer registers per thread. Lower register usage can increase occupancy and improve scheduling flexibility.
• Improved latency hiding: Asynchronous copies allow memory transactions to remain in flight while warps continue executing arithmetic instructions.
• Reduced scoreboard stalls: Traditional load instructions introduce dependencies that prevent subsequent instructions from executing until the data arrives. The asynchronous model delays synchronization until the data is actually required.

Async copies hide latency behind useful work; they do not remove the latency itself.

The effectiveness of the approach therefore depends heavily on the amount of computation available between issuing a copy and consuming the data.

Choosing Pipeline Depth

The most important design decision is often the number of pipeline stages.

Double Buffering

The simplest configuration uses two stages:

Compute Tile N
Load Tile N+1

This approach is easy to implement and requires relatively little shared memory.

For many kernels, double buffering is sufficient because the computation performed on a tile already covers most of the memory latency.

Triple Buffering

Adding a third stage increases the distance between producer and consumer:

Compute Tile N
Load Tile N+1
Load Tile N+2

This can be beneficial when:

• Memory latency is high
• Tile compute is short relative to memory latency
• Global memory bandwidth is not saturated

Triple buffering is common in high-performance GEMM implementations. The benchmark below uses three stages for that reason.

Deeper Stages and Occupancy

Each stage multiplies shared-memory use: for tile size T bytes and depth N, $$\text{Shared Memory} = N \times T$$ A kernel that previously held four blocks per SM may drop to two. Past some depth, occupancy loss outweighs latency hiding. Pick a stage count suited to the workload.

Why Pipelines Sometimes Fail to Improve Performance

Async copies do not automatically speed up a kernel. Benefits appear only when latency is exposed and disappear in two opposite regimes.

Insufficient Computation

Consider a kernel where tile computation only requires a handful of arithmetic instructions.

Copy
Wait
Compute
Copy
Wait
Compute

The pipeline never becomes fully utilized because there is not enough work available to hide memory latency.

Compute Too Fast, or Occupancy Already Hides Latency

At the other extreme, tensor-core GEMM finishes each tile so quickly that staging the next A/B tiles becomes the bottleneck, even at high occupancy. cuBLAS and CUTLASS pipeline global-to-shared loads in exactly this regime.

For bandwidth-bound streaming at full occupancy, software pipelining rarely helps. The kernel is already DRAM-limited and the hardware scheduler is already switching warps.

Synchronization Placement

Issuing memcpy_async and immediately calling consumer_wait before compute eliminates overlap. The copy must be issued far enough ahead of consumption for useful work to run in between.

Warp Specialization and TMA

Warp specialization pushes overlap further by assigning fixed roles within a block, typically one warpgroup as producer (TMA or cp.async loads) and another as consumer (WMMA or WGMMA). Producer and consumer warpgroups advance through the pipeline concurrently instead of alternating copy and compute in the same warps. That separation is what lets Hopper GEMM kernels sustain both high memory throughput and high tensor-core utilization at the same time.

On Hopper, the Tensor Memory Accelerator (TMA) extends the same pipelining idea to bulk transfers: a single thread issues cp.async.bulk.tensor to move an entire tile from global memory into shared memory via a tensor map descriptor, coordinated with mbarriers rather than per-thread cp.async loops. Production library kernels rely on TMA for this staging layer because it cuts register pressure and keeps the copy path off the execution units that run matrix math.

An Experiment: Pipelining a Tensor-Core GEMM

To put these ideas on concrete numbers, I benchmarked a tiled FP16 GEMM (C = A · B, FP32 accumulate) on an H100 against a synchronous baseline. The block tile is 128×128, K is streamed in chunks of 16, and A/B tiles are staged through shared memory with 128-bit loads. The pipelined variant uses cuda::pipeline with three stages. Results and kernels are in tlog-cuda-pipeline.

A three-stage cuda::pipeline delivers a consistent ~1.28–1.31× over synchronous staging. The pipelined kernel is verified against the baseline (cross-checked against a CPU reference at small N). These are teaching kernels: they omit register double-buffering, swizzling, and epilogue tricks a production GEMM uses, so absolute TFLOP/s (~80–85) sits well below cuBLAS. The pipelining effect is still isolated.

Conclusion

CUDA pipelines hide global memory latency by keeping copies in flight while compute runs when that latency is exposed. On an H100 WMMA GEMM, three stages of cuda::pipeline gave a reliable ~1.3× over synchronous staging. gain.

References

1. David Blum, tlog-cuda-pipeline: Pipelining a Tensor-Core GEMM. github.com/davencyw/tlog-cuda-pipeline
2. NVIDIA Developer Blog, Controlling Data Movement to Boost Performance on Ampere Architecture. developer.nvidia.com/blog/controlling-data-movement-to-boost-performance-on-ampere-architecture/
3. NVIDIA, Ampere GPU Architecture Tuning Guide. docs.nvidia.com/cuda/ampere-tuning-guide/
4. NVIDIA, Hopper Architecture In-Depth. developer.nvidia.com/blog/nvidia-hopper-architecture-in-depth/
5. NVIDIA, PTX ISA Documentation (cp.async). docs.nvidia.com/cuda/parallel-thread-execution/
6. NVIDIA CUTLASS Team, Efficient GEMM. github.com/NVIDIA/cutlass/blob/main/media/docs/cpp/efficient_gemm.md
7. Colfax Research, CUTLASS Tutorial: Design of a GEMM Kernel. research.colfax-intl.com/cutlass-tutorial-design-of-a-gemm-kernel/
8. DeepWiki, FMHA Pipeline Coordination and Warp Specialization (CUTLASS). deepwiki.com/NVIDIA/cutlass/6.4-fmha-pipeline-coordination-and-warp-specialization
9. DeepWiki, CUTLASS SM90 MMA Abstractions. deepwiki.com/NVIDIA/cutlass/7.1-mma-abstractions
10. DeepWiki, SM90 TMA Warp-Specialized GEMM (CUTLASS). deepwiki.com/NVIDIA/cutlass/4.2-sm90-tma-warp-specialized-gemm