Asynchronous copy and double buffering

Ampere-class hardware can accelerate appropriate global-to-shared asynchronous copies without the ordinary intermediate per-thread register path. The underlying cp.async family supports documented sizes and alignment requirements, including four-, eight-, and sixteen-byte forms. High-level copy APIs still need suitable trivially copyable types, addresses, sizes, and target support to use their accelerated path. An asynchronous name in source code does not establish which instruction the compiler emitted or whether useful independent work overlaps its execution. Inspect the target code and capability requirements before attributing a speedup to a mechanism.

Hopper adds TMA for supported bulk and multidimensional tensor transfers, using tensor descriptions and architecture-specific layout, alignment, and synchronization rules. It is not merely a larger spelling of every Ampere copy instruction. A kernel port must account for descriptor preparation, valid tile boundaries, and the appropriate completion barrier. For an AI activation tile, preserving the logical tensor mapping matters as much as moving the bytes quickly; a fast copy of a wrongly described stride produces a fast incorrect computation.

A producer acquires a reusable slot, starts its copy, and publishes completion through the mechanism’s required synchronization. A consumer waits for that completion before reading the tile, performs its computation, and releases the slot only after all relevant readers finish. The producer must not overwrite the slot merely because it issued the next copy. Barrier participation and visibility requirements depend on the API and scope. A block-wide barrier placed at a convenient line is not automatically a substitute for the copy operation’s completion protocol.

Double buffering doubles the tile-storage footprint when each stage holds a full tile. That shared-memory cost can reduce resident blocks, while the asynchronous copy path can change register use. More pipeline stages may tolerate longer movement latency but consume additional resources and complicate producer-consumer coordination. Tail tiles, early exits, and divergent participation need particular care because they can break an otherwise regular steady-state schedule. Validate numerical results and synchronization before comparing timing; an occasional stale tile is a correctness failure even when average throughput improves.

Compare a correct sequential baseline with the pipelined kernel on the same GPU, tensor shapes, and numerical precision. Record shared-memory allocation, registers, occupancy constraints, instruction selection, and completion dependencies. Small tile counts spend more time filling and draining; large counts can approach the slower stage’s rate only if resource contention does not invalidate the independent-stage assumption. A timeline or profiler can show whether movement overlaps arithmetic. Merely observing a shorter host submission interval does not demonstrate shorter device execution.

The queue lab offers only an analogy for outstanding tiles. Reset, set arrival rate to 3 and service rate to 2, predict the backlog after one second, then Step twice. Compare a fresh Reset with both arrival rate at 3 and service rate at 4 before two Steps. Its backlog starts at four and changes continuously with fixed rates. It has no copy instructions, buffer slots, barriers, occupancy, or timeline of overlapping stages, so it cannot validate a CUDA pipeline.