Day 74Module 8

draft

Async copies and pipelines

Compile this page's program with -arch=sm_75 and nvcc exits 0. No error, no warning about the missing instruction, and not one cp.async instruction in the PTX; the same source at -arch=sm_80 carries 39 of them (both checked against nvcc 12.6.2 on 2026-09-01).

The API is portable and the instruction is not. A build for the wrong architecture can therefore measure the fallback copy instead.

This lesson needs compute capability 8.0. By the end you can load a GEMM tile with cp.async through a two-stage pipeline, and explain why day 44 stalled at 67.7 percent of cuBLAS.

A copy that skips the registers

Every global-to-shared copy you have written so far uses two instructions: a load from global memory into a register, then a store from the register to shared memory. The thread that issues the load cannot issue the store until the data arrives, so during the fill of a tile, the warp waits through hundreds of cycles of global latency while the FMA units remain idle before a __syncthreads().

cp.async is one instruction that moves 4, 8 or 16 bytes from global memory straight into shared memory without using a register. It exists from compute capability 8.0: "Requires sm_80 or higher." (PTX ISA, cp.async Target ISA Notes, https://docs.nvidia.com/cuda/archive/12.6.2/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async , checked 2026-09-01).

The copy is non-blocking: the warp issues it and continues, then a later operation waits for completion.

A pipeline manages that wait.

cuda::pipeline, from the libcu++ headers day 39 introduced, is a FIFO of stages: producers acquire the head stage, submit cuda::memcpy_async calls into it, and commit it; consumers wait on the tail stage, compute on its data, and release it.

With two stages, the block computes on a tile fetched by the previous iteration while the next fetch is in flight. Hopper's tensor-core pipelines, TMA and warp-specialised GEMMs use the same pattern with larger copies and dedicated hardware. Days 75 to 78 build on it.

The same overlap at a different scope

Day 54 pipelined a 1 GB array through a 64 MB window: two host buffers, cudaMemcpyAsync on two streams, compute on one chunk while the next copies. The order is similar, but the mechanism differs.

There is no copy engine in this lesson; cp.async is issued by the warp itself and the data moves through the SM's own memory pipe. There is no stream and no event; the "async" is asynchronous with respect to the issuing thread, inside one kernel, and the pipeline object is the only thing that can wait on it.

Do not infer the instruction from the API alone. The CUDA 12.6 guide says the memcpy_async APIs "require compute capability 7.0 or higher", and that on "8.0 or higher" they "can benefit from hardware acceleration" (https://docs.nvidia.com/cuda/archive/12.6.2/cuda-c-programming-guide/index.html , section 7.27.1, checked 2026-09-01).

Below 8.0, or above it with the wrong -arch flag, the same source compiles to the register-path copy, correct and unaccelerated. Nothing fails. Your measurement just stops being about cp.async.

Two kernels, one arithmetic

Full program in code/day74-cp-async/pipelined_gemm.cu. It is a smaller form of day 43's tiled matmul: a 64 by 64 block tile, 16-deep K slices, a 4 by 4 micro-tile per thread, 256 threads. Both kernels call one shared computeTile() function, so they execute identical floating-point operations in an identical order, and the program requires their outputs to be bit-identical with memcmp before it times anything.

The fill is the only difference between them, so the fill is the only thing the ratio can measure. Each thread moves one float4 of the A slice and one of the B slice per K step, at the same addresses in both kernels.

The baseline fills through registers and barriers:

    for (int kt = 0; kt < size / kBlockK; ++kt) {
        const int kBase = kt * kBlockK;
        *reinterpret_cast<float4*>(&tileA[rowA * kBlockK + colA]) =
            *reinterpret_cast<const float4*>(
                &a[(rowBase + rowA) * size + kBase + colA]);
        *reinterpret_cast<float4*>(&tileB[rowB * kBlockN + colB]) =
            *reinterpret_cast<const float4*>(
                &b[(kBase + rowB) * size + colBase + colB]);
        __syncthreads();

        computeTile(tileA, tileB, threadRow, threadCol, acc);
        __syncthreads();
    }

The pipelined kernel doubles the tile storage and adds the pipeline's shared state, which cuda::make_pipeline initializes collectively:

    // cp.async's 16-byte form needs 16-byte-aligned shared addresses; a
    // bare float array only promises 4.
    __shared__ alignas(16) float tileA[kStages][kBlockM * kBlockK];
    __shared__ alignas(16) float tileB[kStages][kBlockK * kBlockN];
    __shared__
    cuda::pipeline_shared_state<cuda::thread_scope::thread_scope_block, kStages>
        state;

    auto block = cooperative_groups::this_thread_block();
    auto pipe = cuda::make_pipeline(block, &state);

Its main loop keeps kStages fetches ahead of the compute. The inner loop tops up the pipeline; producer_acquire blocks only when every stage slot is full, which is what bounds the lookahead:

    const int tiles = size / kBlockK;
    for (int compute = 0, fetch = 0; compute < tiles; ++compute) {
        for (; fetch < tiles && fetch < compute + kStages; ++fetch) {
            pipe.producer_acquire();
            const int s = fetch % kStages;
            const int kBase = fetch * kBlockK;
            cuda::memcpy_async(&tileA[s][rowA * kBlockK + colA],
                               &a[(rowBase + rowA) * size + kBase + colA],
                               cuda::aligned_size_t<16>(sizeof(float4)), pipe);
            cuda::memcpy_async(&tileB[s][rowB * kBlockN + colB],
                               &b[(kBase + rowB) * size + colBase + colB],
                               cuda::aligned_size_t<16>(sizeof(float4)), pipe);
            pipe.producer_commit();
        }

        pipe.consumer_wait();
        computeTile(tileA[compute % kStages], tileB[compute % kStages],
                    threadRow, threadCol, acc);
        pipe.consumer_release();
    }

cuda::aligned_size_t<16> asserts two facts: it tells the library both pointers are 16-byte aligned and the size is a multiple of 16, which is what lets each call lower to a single 16-byte cp.async. The guide states the consequence: "If the proof is incorrect, the behavior is undefined" (section 7.27.6.1, checked 2026-09-01). In this program the constants keep the promise, and the static_asserts refuse shapes that would break it.

Note. Neither kernel guards its bounds; every size is a whole number of tiles in M, N and K, enforced at compile time. Real libraries pay for ragged edges with a separate epilogue path. Day 44 made the same trade for the same reason: the guard would sit in the inner loop, the one place this comparison cannot afford it.

Results

Not yet run. Verification is blocked on hardware: the T4 node is CC 7.5, and this program's capability gate exits there by design. One thing has been graded, at the toolkit version the T4 node runs.

At -arch=sm_80 nvcc 12.6.2 exits 0 with a single warning about the pipeline state's initialization, and the PTX carries 39 cp.async lines, all of them inside matmulTilePiped: 33 cp.async.cg.shared.global, 3 cp.async.mbarrier.arrive.shared.b64, 3 cp.async.wait_all. matmulTileSync has none, at either arch.

That is a compile, not a run: the transcript is code/day74-cp-async/evidence/compile-2026-09-01.txt and it says in its first line that no GPU touched it. Every number in the table below is still missing, and the run that fills it happens on a CC 8.0 or newer card.

N	sync ms	piped ms	GFLOP/s each	piped/sync
512	?	?	?	?
1024	?	?	?	?
2048	?	?	?	?

The first sm_80 run must check these predictions. Day 44 predicted 70 to 80 percent and measured 67.7 percent, so each range here can also fail.

The piped kernel wins, by 5 to 25 percent at N = 2048 on an A100-class card, so piped/sync lands between 0.75 and 0.95. The fill is a minority of each iteration's work, so the gain is capped by its share.

At or above 1.0, the pipeline's own bookkeeping used the expected gain. Below 0.75, something other than the fill changed.
memcmp passes at every size. The kernels share their arithmetic, so scheduling is the only thing the pipeline may change. A mismatch is a fill bug, and the program says so and exits nonzero.
cuobjdump -sass shows LDGSTS in the piped kernel. The open half of this bet is the sync kernel: ptxas is allowed to fuse an LDG-STS pair into LDGSTS on sm_80 on its own, and if it has, the timing gap narrows and this page gains a paragraph explaining why. The SASS capture settles it either way.

Absolute timings differ by GPU. Compare the ratio, which shows whether the fill overlapped the FMAs, and the SASS, which identifies the copy instruction that ran.

Run it yourself

Run the binary on a GPU with compute capability 8.0 or newer. The -arch=sm_80 build embeds PTX that can JIT-compile for later capabilities. The build line and the compile-only checks for older cards are in the README.

For ways to access compatible hardware, see /setup/learn-cuda-without-a-gpu. Cost notes for the course's test runs remain in FACT-SHEET.md section 4.

Exercise

kStages is one constant. Build at 2, 3 and 4 stages, run all three at N = 2048, and explain the shape of the three piped/sync ratios in two sentences.

Time: 30 to 45 minutes. Submit: the three ratios, the shared memory per block at each stage count, and one sentence naming what a deeper pipeline buys and what it spends.

Check: report ratios because absolute times differ by card. Correctness still gates every build: a stage-indexing mistake fails the memcmp against the sync kernel, and the program prints the size and exits nonzero before any timing.

The static_asserts accept exactly 2, 3 and 4 stages, and that ceiling is the exercise's, not the card's: at 5 stages ptxas still reports 41048 bytes of shared memory per block, inside the 48 KiB static limit. Six is the first depth the limit itself refuses.

Hint 1

Each extra stage is one more tile fetch allowed in flight, and one more copy of both tiles resident in shared memory. What is the latency the lookahead exists to cover, and how many stages' worth of compute already covers it?

Hint 2

One stage of tiles is 8 KiB, and ptxas reports the block's shared-memory total climbing 16424, 24632, 32840 bytes across 2, 3 and 4 stages (the compile transcript in the code directory). Check what those do to resident blocks per SM on your card before judging the pipeline, and remember consumer_wait only ever waits on the oldest stage: if that copy finished long ago, the deeper queue changed nothing.

Solution

Expect 3 stages to match or slightly beat 2, and 4 to be flat or worse. Two stages already overlap each fetch with a full tile's compute, 256 FMAs per thread; if that compute time exceeds the global latency of the next fetch, the copy is done before consumer_wait asks, and extra stages just queue completed work. What the deeper pipeline spends is real either way: 8 KiB of shared memory per stage, which at four stages can cut resident blocks per SM and cost the occupancy that hides every other stall.

Stage depth can cover more latency but uses more shared memory. Use the smallest depth that hides the fetch.

Pitfalls

You built with -arch=sm_75 and everything works, slower than promised. The wrong arch flag does not fail; nvcc 12.6.2 exits 0 and emits zero cp.async instructions, so the program runs the register fallback and measures it (verified compile-only, 2026-09-01). Check the PTX or SASS before trusting a pipeline number, the way day 46 reads it.

The sm_80 binary refuses to run on an older card. The string is no kernel image is available for execution on the device: compute_80 PTX JIT-compiles forward, never backward. There is no fix except CC 8.0 hardware; this program checks the capability first so the message names the card instead. Day 69 covers shipping one binary for both.

The kernel hangs. Every producer_acquire needs its producer_commit, every consumer_wait its consumer_release, and all threads of the block must make each call: the pipeline is a block-scope object and a partial commit strands the whole block. Diagnosing it is day 65's job; preventing it is loop structure, which is why the fetch top-up and the compute live in one loop here.

You waited with __syncthreads() instead of consumer_wait(). A barrier orders threads; it knows nothing about copies in flight. The PTX ISA is explicit that no other synchronization "can be used to guarantee the completion of the asynchronous copy operations" (cp.async section, checked 2026-09-01). Reading a tile after a barrier but before the wait is a race that may pass when other work is low.

Your aligned_size_t promise is false and the output is garbage or worse. The alignment argument is a proof you assert, and "If the proof is incorrect, the behavior is undefined." A size that stops being a multiple of 16 bytes, or a shared array without alignas(16), breaks it silently. This program's static_asserts pin both; carry that habit.

You expected 2x and got 8 percent. Overlap can only hide the smaller of copy and compute. In this kernel the FMAs dominate each iteration, so hiding the whole fill moves the total by the fill's share and no more.

Latency hiding reduces wait time; it does not multiply throughput. The roofline from day 30 shows whether memory or compute limits the kernel.

Go deeper

CUDA C++ Programming Guide 12.6, sections 7.27 "Asynchronous Data Copies" and 7.28 "Asynchronous Data Copies using cuda::pipeline": https://docs.nvidia.com/cuda/archive/12.6.2/cuda-c-programming-guide/index.html (checked 2026-09-01)
PTX ISA, cp.async and its commit and wait groups: https://docs.nvidia.com/cuda/archive/12.6.2/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cp-async (checked 2026-09-01)
cuda-samples, cpp/3_CUDA_Features/globalToShmemAsyncCopy: https://github.com/NVIDIA/cuda-samples/tree/master/cpp/3_CUDA_Features/globalToShmemAsyncCopy (checked 2026-09-01)
CUTLASS, "Efficient GEMM in CUDA", the pipelining section, where this two-stage shape is the production default: https://github.com/NVIDIA/cutlass/blob/main/media/docs/cpp/efficient_gemm.md (checked 2026-08-30)

Day 75 covers the primitive underneath consumer_wait: the asynchronous barrier, which lets a thread announce arrival and keep working before it waits, and which is how producer and consumer stop being the same threads. Day 76 replaces the per-thread cp.async loop with one TMA descriptor that moves the whole tile, and day 80's capstone stretch goal is exactly this page's pipeline feeding mma.sync instead of FMAs.