100 DAYS OF GPU Course preview
Day 54Module 6
in-technical-review

Double buffering

Here is a program with nothing wrong in it. It takes a 1 GiB array, copies a 64 MiB chunk to the device, runs the kernel on it, copies the result back, and repeats sixteen times. Every call is checked, the host memory is pinned, and the answer is right.

Its time is the copy in, plus the compute, plus the copy out, because that is the order the source says and the order one queue runs.

The GPU has copy engines separate from its SMs, and they can work at the same time. This program keeps all but one type of unit idle at every moment.

How many copy engines, and whether the two directions get one each, is what the program's asyncEngineCount line reports and what prediction 4 bets on. By the end of this page you will run the same sixteen chunks through two buffers and two streams, watch the three rows interleave in Nsight Systems, and know why the total falls to the slowest of the three stages instead of their sum.

The slowest stage sets the minimum time

Chunking splits one big job into repeated work. Each chunk passes through three stages: host to device across PCIe, the kernel on the SMs, device to host.

Copies and kernels run on different resources: a copy costs no SM cycles and a kernel costs no PCIe bytes. Whether the two copy directions also get separate engines is prediction 4's bet, not a given.

What serializes them in the naive program is not the hardware but the queue. A stream is an ordered queue, and day 51 showed that order inside one stream is absolute: chunk 3's copy in cannot start until chunk 2's copy out has finished, even though they use different engines and different buffers.

Double buffering removes the false dependency by adding a second device buffer pair. Chunk k+1 then has space while chunk k is still being read.

Even chunks own pair 0 and stream 0, odd chunks own pair 1 and stream 1. Now the H2D engine can pull chunk 1 in while the SMs process chunk 0.

Once both buffers are in use, each hardware unit can work on a different chunk. The timeline then shows operations from the two streams overlapping.

The steady state has a simple cost model. With N chunks, the total includes the first chunk's copy in, N steps at the pace of the slowest stage, and the last chunk's copy out.

For large N, the main cost is N times the slowest stage because the other stages overlap it. The naive program pays the sum of the three stages; the pipeline pays their maximum. The difference between sum and maximum is the value of this pattern, and it is a number this page's program prints.

Two streams do not mean twice as fast

Concurrency does not divide the time by the number of streams. Overlap runs the shorter stages at the same time as the slowest one; it does not shorten that stage.

If the copy in is three quarters of the total, the perfect pipeline still costs the copy in, and the speedup tops out well short of two, no matter how many streams you add. Day 51 gives the fix for a slow stage: move less data, or keep it resident.

The first and last chunks also add serial work. No earlier work can overlap the first chunk's copy in, and no later work can overlap the last chunk's copy out.

Those two operations explain why a four-chunk pipeline is worse than a sixteen-chunk one: the edges are the same size, but they are half of four and an eighth of sixteen. The exercise makes you measure exactly that.

One schedule moves, nothing else

Full program in code/day54-double-buffer/double_buffer.cu. Four commitments make the schedule the only variable.

Every variant does identical work. The same 2 GiB crosses PCIe, the same sixteen kernels run; only the schedule moves. The program also times three component passes (all copies in, all kernels, all copies out) so the sum and the max are measured, not inferred.

Pinned host memory on both ends. Both 1 GiB host buffers come from cudaMallocHost, because an asynchronous cudaMemcpyAsync from pageable memory degrades to a staged copy and the overlap silently disappears. Day 53 measured that difference; this program spends it.

Buffers belong to streams, and stream order is the only lock. Chunk k+2 reuses chunk k's buffers, and both sit in the same stream, so the reuse cannot start early. The schedule needs no events, syncs, or flags.

To test that claim instead of assuming it, correctness runs before any timing: the naive output is checked against a double-precision closed form, and the pipelined pass must then reproduce it bit for bit. A buffer race here is a failing exit with a chunk number, not noise in a table.

CUDA events time it, from the legacy stream. The two worker streams are created blocking, so an event recorded on the legacy default stream is a barrier: start fires before any timed work is queued, stop completes when both streams have drained, and the host waits at cudaEventSynchronize(stop) before reading the clock. Each timed section is wrapped in an NVTX range, day 41's habit.

Here is the naive schedule, one stream, one buffer pair:

static void passNaive(const Ctx& c) {
    for (int k = 0; k < kChunks; ++k) {
        const size_t off = static_cast<size_t>(k) * kChunkElems;
        CUDA_CHECK(cudaMemcpyAsync(c.d_in[0], c.h_in + off, kChunkBytes,
                                   cudaMemcpyHostToDevice, c.stream[0]));
        launchChunk(c.d_in[0], c.d_out[0], c.stream[0]);
        CUDA_CHECK(cudaMemcpyAsync(c.h_out + off, c.d_out[0], kChunkBytes,
                                   cudaMemcpyDeviceToHost, c.stream[0]));
    }
}

And the pipeline. The diff is three characters: [0] becomes [s].

static void passPipelined(const Ctx& c) {
    for (int k = 0; k < kChunks; ++k) {
        const int s = k & 1;
        const size_t off = static_cast<size_t>(k) * kChunkElems;
        CUDA_CHECK(cudaMemcpyAsync(c.d_in[s], c.h_in + off, kChunkBytes,
                                   cudaMemcpyHostToDevice, c.stream[s]));
        launchChunk(c.d_in[s], c.d_out[s], c.stream[s]);
        CUDA_CHECK(cudaMemcpyAsync(c.h_out + off, c.d_out[s], kChunkBytes,
                                   cudaMemcpyDeviceToHost, c.stream[s]));
    }
}

Note. The kernel runs 512 chained FMAs per element, which exists to give the compute stage enough work to measure. The test GPU holds this whole GiB. The same pattern also works when the array exceeds device memory or arrives in chunks from another source.

Results

Re-verified 2026-09-02 on the same Tesla T4 with driver 580.173.02, CUDA 13.0 (V13.0.88), and the same three reported copy engines. Naive moved from 212.378 to 211.246 ms, double-buffered from 116.889 to 115.465 ms, and the speedup held at 1.83x.

Both transcripts and the fresh CUDA 13 report are listed in front matter. The published CUDA 12.6 report and its CSV tables ship at content/profiles/double-buffering/, so the exercise is doable without a GPU: day54-pipeline.nsys-rep under -t cuda,nvtx, with nvtx_sum, cuda_gpu_mem_time_sum and cuda_api_sum exports beside it.

Pass Mean (ms) GB/s
h2d only 87.809 12.23
kernel only 43.700
d2h only 82.444 13.02
naive bound (sum) 213.953
pipeline floor (max) 87.809
measured naive 211.246
measured double buffer 115.465

What the arithmetic promises

  1. The measured naive pass lands within 10 percent of the sum of the three components. One stream serializes exactly; if naive comes in clearly under the sum, something in it overlapped and my model of the legacy queue is wrong.
  2. The pipelined pass lands between the slowest component and that floor plus one chunk's copy in and one chunk's copy out. The edges are the only serial residue this schedule cannot hide.
  3. The copies are the two slowest stages, each several times the kernel pass. Day 9 paid 45.822 ms to move three 64 MiB buffers and run one trivial kernel, and this node's PCIe link has not gotten faster since.
  4. The pipelined pass beats naive by at least 1.6x. This result needs two or more copy engines (the program prints asyncEngineCount), so the copy in of one chunk overlaps the copy out of another. If the card reports one engine, this prediction dies, the two copy stages share a lane, and the saving shrinks to roughly the kernel stage.

What the run said

  1. Held. Naive measured 211.246 ms against a 213.953 ms sum of components, 2 percent under. One stream serializes.
  2. Failed, which exposes the model's error. The prediction allowed the pipeline the slowest stage plus two chunk edges, 87.809 plus about 10.6 ms. Measured is 115.465 ms, 27.656 ms over the floor. The model was wrong about what two streams can hide: each stream still runs its own eight chunks serially (copy in, kernel, copy out, in stream order), so each stream's chain costs half of 213.953, about 107.0 ms, and the card can only overlap one stream's copies with the other stream's work. 115.465 sits just above that 107.0 bound, not above the 87.809 one. Hiding all three stages at once needs at least three buffers in flight, not two; the exercise's window sweep is where to feel that.
  3. Half held. The copies are the two slowest stages (87.809 and 82.444 ms against 43.700), but at about twice the kernel pass each, not "several times". The kernel is heavier relative to this link than day 9's trivial one was.
  4. Held. The card reports 3 copy engines and the pipeline beat naive by 1.83x, past the 1.6x bar.

The useful result across systems is which stage is slowest. Measure the stages before pipelining anything; if one of them is most of the total, overlap can only refund the rest.

Run it yourself

Use a CUDA GPU with about 3 GiB of free host RAM (2 GiB of it pinned) and 256 MiB of device memory.

nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo -o double_buffer double_buffer.cu
nsys profile -t cuda,nvtx -o profile/day54-pipeline ./double_buffer

There is no Compiler Explorer embed: the program runs 32 full passes over the array (two verify passes, then a warm-up plus five timed runs for each of its five variants) and moves about 40 GiB across PCIe doing it, far past the 20 second execution cap. If you have no GPU, read /setup/learn-cuda-without-a-gpu and do the exercise from the shipped reports; the memory rows of the timeline are the point, and they open in the free Nsight Systems GUI on any laptop.

Exercise

Change the window: set kChunkElems to a 4 MiB chunk, then a 256 MiB chunk, rerun each, and record the pipelined-to-naive ratio all three times. Explain which end of the sweep got worse and why, using the edge argument and the per-chunk overhead argument, one each.

Time: 30 to 45 minutes. Submit: the three ratios and two sentences.

Check: compare ratios rather than milliseconds so results remain useful across GPUs. The 64 MiB window should put pipelined at or under 70 percent of naive. The cost model says the 256 MiB window (four chunks) gives the worst ratio of the three: if yours does not, your runs disagree with the model and one of them needs a timeline capture.

The shipped day54-pipeline_nvtx_sum.csv gives the reference ratios if you cannot run it.

Hint 1

At the moment the very first chunk is crossing PCIe, what else is the card doing? Same question for the moment the very last chunk crosses back. How much of the whole run are those two moments when there are four chunks, and when there are 256?

Hint 2

The other end of the sweep is not free either. Every chunk costs three API calls whatever its size; 256 chunks cost 768 enqueues plus a launch overhead per kernel. Compare cuda_api_sum between your 4 MiB run and the 64 MiB run: does the count grow faster than the time?

Solution

Total pipeline time is roughly fill plus N-1 steps of the slowest stage plus drain. Shrinking chunks reduces the first and last serial operations, which is why the model says 4 MiB should beat or tie 64 MiB: the per-chunk overheads are launch-scale (day 48 measured 2.668 microseconds per launch on this card) while a chunk's copy time is milliseconds-scale, so 256 of them should still cost almost nothing. The run prices both sides of that bet.

Growing chunks to 256 MiB leaves only four, so fill and drain take a large share of the total and the ratio moves toward naive.

Estimate pipeline time from the slowest stage times the chunk count, plus the first and last serial operations. More chunks spread those two operations across more work.

Pitfalls

Your streams overlap nothing, and the timeline shows every copy on one row. The host buffer is pageable.

cudaMemcpyAsync needs pinned memory to be asynchronous; from pageable memory it stages through a pinned bounce buffer and behaves like the blocking copy. No error, no warning. Day 53 is the fix.

One launch in the middle uses the default stream and stops the overlap. A kernel launched with <<<blocks, threads>>> and no stream argument lands in the legacy default stream, which serializes against both worker streams: everything before it must drain, everything after it must wait. One forgotten fourth argument turns the pipeline back into a chain at that point.

A cudaDeviceSynchronize inside the chunk loop. Usually left over from debugging. It drains both streams once per chunk, so no chunk ever overlaps its neighbour and the pipelined pass times the same as naive. Sync once, after the loop.

You timed the enqueue, not the work. A host clock around the pipelined loop measures sixteen asynchronous enqueues returning immediately, a fraction of a millisecond, and the conclusion is a thousand-fold speedup. Day 9 is the same bug in its simplest form. Synchronise before reading any clock, or use events as this program does.

Two buffers, three streams. The buffer count, not the stream count, sets the pipeline depth. With more streams than buffer pairs, a chunk can be scheduled into a buffer whose previous tenant has not been read out, and stream order no longer protects it because a different stream reuses the buffer.

The result is wrong output without an API error, which is why this program bit-compares the two schedules. Cross-stream protection needs day 52's events.

Go deeper

Next

Day 55 removes the other serializer hiding in loops like this one: a cudaMalloc per iteration, which blocks the whole device where cudaMallocAsync joins the queue. Then day 56 takes the schedule you built by hand today and captures it as a graph, so the sixteen-chunk enqueue loop becomes one launch.