Day 60Module 6

in-technical-review

Capstone 3: a real-time frame pipeline

Your task is to deliver 900 synthetic 1920x1080 frames at 60 frames per second through host generation, two copies, and a three-kernel CUDA graph. The result passes only when frames 0, 5, 7, and 899 match the CPU reference, the drop counter is zero, and the Nsight Systems trace shows the intended overlap in the burst phase.

This capstone joins the scheduling skills from days 51 through 59. You will predict the limit, trace one frame through the dependencies, run the full pipeline, and compare its sustained rate with the limit printed by the program.

Retrieve one rule

Day 54 compared a serial three-stage schedule with a double-buffered schedule. Before reading on, write the two expressions for their steady-state cost.

Check your answer

The serial cost is the sum of the stage times. With enough chunks and valid overlap, the pipelined cost approaches the largest stage time.

Day 53 adds one required condition: host buffers used by cudaMemcpyAsync must be pinned so a copy engine can use them without pageable staging.

Predict the schedule

One frame has four stages: host generation, a 2,073,600-byte H2D copy across PCIe, a blur-Sobel-threshold chain, and a D2H copy. Cadence is the target interval between frames; at 60 frames per second, it is 16.667 ms, or about 16.7 ms.

Predict whether the sum or maximum of the four stage times sets the full pipeline's maximum rate. Then predict whether the host generator, a copy, or the kernel chain will be the largest stage on your machine.

Make three more predictions before you run the program:

Will the synchronous baseline take at least 1.5 times as long per frame as the burst pipeline?
Which pair should overlap in the burst trace: frame n's chain with frame n+1's upload, or two operations from the same stream?
Will the paced trace show dense overlap or mostly idle time at 60 fps?

Keep these answers. The Results section comes after the exercise so the measurements cannot replace your prediction.

Day 41 taught you how to read the copy and kernel rows that will test the second and third predictions.

A small performance model

Let G, U, K, and D be the per-frame times for generation, upload, the kernel chain, and download. A serial frame costs T_serial = G + U + K + D.

In steady state, different resources can work on different frames. If the dependencies permit overlap, the lower bound is T_pipeline = max(G, U, K, D), and the rate ceiling is 1000 / T_pipeline frames per second.

The device-only burst uses pregenerated frames, so its bound is max(U, K, D). The paced run generates each frame, so its bound includes G.

This maximum is a bound, not a measured rate. Queue gaps from launch overhead, a cudaDeviceSynchronize inside the loop, or a false dependency can keep the pipeline below it. Day 9 showed why the work around a kernel belongs in the model.

A serial schedule pays the sum of generation, upload, kernels, and download. A pipeline overlaps stages from different frames, so its steady-state cost approaches the longest stage and must fit within 16.667 milliseconds at 60 frames per second.

Work one frame through the schedule

The program uses three streams: upload, compute, and download. A CUDA event records completion in one stream, and cudaStreamWaitEvent makes another stream wait without blocking the host. Each transfer uses cudaMemcpyAsync.

Trace frame 2, which reuses device buffer 0 after frame 0:

The upload stream waits for graphDone[0], then copies frame 2 into d_in[0]. This wait prevents the copy from overwriting frame 0's input while its graph still reads it.
The compute stream waits for frame 2's copyDone[0] and frame 0's frameDone[0]. The first wait protects the new input, and the second protects the output buffer that the old download still reads.
The compute stream launches graph 0 and records graphDone[0]. The download stream waits for that event before it reads the edge buffer.
The H2D copy records slotCopied[2] after it has read the pinned ring slot. The host waits for this event before it lets the producer reuse that slot; it does not wait for the kernels or D2H copy.

Frames 0 and 1 use the two device buffer sets. Once both sets are active, frame n+1 can upload while frame n computes and frame n-1 downloads.

The full program is in code/day60-capstone-3/, with the source in frame_pipeline.cu. The three kernels are captured once for each buffer parity using the CUDA graph method from day 56, so 900 frames use two graph executables rather than graph updates. Day 57's graph update machinery is not needed because the two captured pointer sets do not change.

static cudaGraphExec_t captureChain(cudaStream_t stream, unsigned char* in,
                                    unsigned char* blur, unsigned char* grad,
                                    unsigned char* edges) {
    const dim3 block(kBlockDimX, kBlockDimY);
    const dim3 grid(kGridX, kGridY);
    cudaGraph_t graph = nullptr;
    CUDA_CHECK(cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal));
    boxBlur3<<<grid, block, 0, stream>>>(in, blur, kRows, kCols);
    CUDA_CHECK(cudaGetLastError());
    sobelMag<<<grid, block, 0, stream>>>(blur, grad, kRows, kCols);
    CUDA_CHECK(cudaGetLastError());
    thresholdBinary<<<grid, block, 0, stream>>>(grad, edges, kRows, kCols);
    CUDA_CHECK(cudaGetLastError());
    cudaGraphExec_t exec = nullptr;
    CUDA_CHECK(cudaStreamEndCapture(stream, &graph));
    CUDA_CHECK(cudaGraphInstantiate(&exec, graph, 0));
    CUDA_CHECK(cudaGraphDestroy(graph));
    return exec;
}

Here is the per-frame issue path. Match each wait to the reason above.

static void issueFrame(const Pipeline& p, int n) {
    const int buf = n & 1;
    const int slot = n % kRingSlots;

    nvtxRangePushA("upload");
    CUDA_CHECK(cudaStreamWaitEvent(p.copyStream, p.graphDone[buf], 0));
    CUDA_CHECK(cudaMemcpyAsync(p.d_in[buf], p.h_ring + slot * kFrameBytes,
                               kFrameBytes, cudaMemcpyHostToDevice,
                               p.copyStream));
    CUDA_CHECK(cudaEventRecord(p.slotCopied[slot], p.copyStream));
    CUDA_CHECK(cudaEventRecord(p.copyDone[buf], p.copyStream));
    nvtxRangePop();

    nvtxRangePushA("chain");
    CUDA_CHECK(cudaStreamWaitEvent(p.computeStream, p.copyDone[buf], 0));
    CUDA_CHECK(cudaStreamWaitEvent(p.computeStream, p.frameDone[buf], 0));
    CUDA_CHECK(cudaGraphLaunch(p.chain[buf], p.computeStream));
    CUDA_CHECK(cudaEventRecord(p.graphDone[buf], p.computeStream));
    nvtxRangePop();

    nvtxRangePushA("download");
    CUDA_CHECK(cudaStreamWaitEvent(p.downloadStream, p.graphDone[buf], 0));
    CUDA_CHECK(cudaMemcpyAsync(p.h_out + buf * kFrameBytes, p.d_edges[buf],
                               kFrameBytes, cudaMemcpyDeviceToHost,
                               p.downloadStream));
    CUDA_CHECK(cudaEventRecord(p.frameDone[buf], p.downloadStream));
    nvtxRangePop();

    CUDA_CHECK(cudaEventSynchronize(p.slotCopied[slot]));
    gReleased.store(n + 1, std::memory_order_release);
}

Build the correctness check

A std::thread, using the pattern from day 59, writes each frame into a four-slot pinned ring. The frame is a pure function of its index, and the CPU and GPU share the same __host__ __device__ integer arithmetic, so the program compares their outputs with memcmp.

Frames 5 and 7 force reuse beyond the two device buffers and four ring slots. A missing event then produces a wrong byte instead of remaining hidden on a fast run.

The paced producer counts a frame as dropped when its ring slot remains occupied at that frame's tick. A camera would overwrite the frame, but this program counts and still delivers it so frame 899 remains checkable. The program gate is dropped == 0.

CUDA events time each device component on its worker stream, with warm-up inside the timer. The host monotonic clock measures generation and the full pipeline because those spans cross one host thread and three streams; each read occurs after the synchronization that settles its work.

The source names this host-clock rule in wallMs, and captureChain keeps device synchronization outside graph capture. The producer's generate range contains only host work.

Check the model

Answer these before adding more code:

Why does the compute stream wait for frameDone[buf] as well as copyDone[buf]?
Why can the producer reuse a ring slot after slotCopied[slot], before the graph and download finish?
Which trace range proves capacity, burst or paced?

Answers

The first wait protects the reused edge buffer from an unfinished download. The ring slot is safe once H2D has read it because later stages use device memory. The unpaced burst range proves capacity; the 60 fps paced range mainly proves cadence and should contain slack.

The three kernels do light per-pixel work. Day 20 measured 3.242 ms for its naive blur and 2.619 ms for its separable blur on a 4096x4096 image, which has about eight times as many pixels as this frame.

Faded practice

First, draw frames 0, 1, and 2 on upload, compute, and download rows. Label the four stream waits for frame 2 with the buffer or data they protect; use the worked example above if needed.

Next, draw frame 3 without looking back. Your schedule must allow frame 3's upload to overlap frame 2's graph while preserving device buffer parity and ring slot ownership.

Finally, remove one event wait from your drawing. Name the first frame whose bytes may be overwritten and which correctness check can expose it.

Run it yourself

Run this on a CUDA-capable GPU. nsys traces as a normal user, so the ERR_NVGPUCTRPERM error from day 42 does not apply.

nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo -o frame_pipeline \
    frame_pipeline.cu
./frame_pipeline
nsys profile --trace cuda,nvtx --force-overwrite true \
    --output day60-pipeline ./frame_pipeline

The program pins host memory, runs a producer thread, and spends 15 seconds in the paced phase. It does not have a Compiler Explorer embed.

Without a CUDA GPU, complete the analysis from the shipped report and CSV exports in content/profiles/frame-pipeline-capstone/.

Independent exercise

Find your card's breaking cadence. Raise kTargetFps in steps (90, 120, 180, 240), rebuilding and rerunning, until the drop counter stops being zero. Then say which line of the components table predicted that number.

Time: 30 to 45 minutes.

Submit your components table, the highest target that still passes, the first target that fails, and one sentence naming the limiting stage.

The correctness gate is exact: the unchanged 60 fps run must report frames 0, 5, 7, and 899 matching the CPU reference, followed by PASS: dropped 0 over 900 frames at 60 fps.

The performance gate is also exact. The highest passing target must report dropped 0, the next tested target must report a nonzero drop count, and 1000 / first_failing_target must be within 25 percent of the largest printed per-frame component.

If it breaks at a much lower rate than the components allow, something serializes. Open the timeline, find the stream that waits where the burst phase shows it should not, and fix that instead of the kernels.

Hint 1

The program already printed the answer before you changed anything. Two ceilings appear in the components block. Which one is the paced run actually up against, and what rate does it correspond to?

Hint 2

In the paced capture, compare the width of the generate range on the producer thread against the kernel row's busy time inside one tick. The pipeline reaches its limit at its slowest stage. Which stage is slowest on your machine?

Worked answer for the reference run

In the reference run, the expected limit is the rate with live generation, set by the host generator. The breaking cadence should track 1000 divided by the printed generate cost, and the copies and kernels still have room when it breaks. The fix that raises the ceiling is therefore not a faster kernel but a cheaper or parallel producer: split synthesis across two host threads feeding the same ring and the breaking cadence should nearly double, while touching the CUDA code not at all.

A pipeline reaches its limit at its slowest stage. Grade the result against its printed ceiling so the same check works on each GPU.

Results

The measurements below come after the prediction and exercise design. They are evidence for this implementation, not a substitute for measuring your own machine.

The program was re-verified on 2026-09-02 on a Tesla T4 with driver 580.173.02, CUDA 13.0 (V13.0.88), and three reported copy engines. Frames 0, 5, 7, and 899 matched the CPU reference, and all 60 fps gates passed.

The CUDA 13 run moved the device ceiling from the CUDA 12.6 run's 0.314 ms H2D copy to the 0.188 ms kernel chain. The host generator remained the largest full-pipeline component.

Component (per frame)	ms
generate (host)	1.547
H2D copy	0.178
kernel chain	0.188
D2H copy	0.163

The printed ceilings: 0.188 ms per frame (kernel chain) is 5306.6 fps device-only; 1.547 ms per frame (generate) is 646.2 fps with live generation. The budget at 60 fps is 16.667 ms per frame.

Phase	ms per frame	Rate	Gate
baseline, synchronous	0.846	1182.7 fps	none
burst, pipelined	0.239	4180.3 fps	none
paced at 60 fps	16.651	60.05 fps	dropped 0 of 900: PASS

The paced run delivered 900 frames with zero drops at 60.05 fps, within one percent of the target. The generator took 1.547 ms, over eight times the 0.188 ms kernel chain, so it set the live-generation ceiling.

The synchronous baseline took 0.846 ms per frame and the burst took 0.239 ms, a 3.54x ratio. The burst reached 0.79 of the device ceiling, and its 0.239 ms result leaves no serial room for a 0.188 ms chain beside a 0.178 ms copy.

The published CUDA 12.6 trace shows the burst memcpy row active under the kernel row while the baseline rows alternate. In the paced range, the three device stages occupy less than one third of each 16.667 ms tick, and the widest per-frame NVTX range is generate on the producer thread.

That CUDA 12.6 paced run sustained 60.06 fps. The CUDA 13 table above reports 60.05 fps after rounding its new transcript.

Both transcripts and the CUDA 13 report are listed in front matter. The published CUDA 12.6 report and CSV exports are in content/profiles/frame-pipeline-capstone/.

Compare each prediction with these results. If your sustained rate is far below your printed ceiling, use the burst timeline to name the queue gap or wait that accounts for the difference.

Pitfalls

cudaMemcpyAsync from pageable memory can fall back to staged, synchronous-looking behavior while the call still succeeds. The CUDA Runtime API synchronization rules describe this behavior; if the trace shows no copy-compute overlap, allocate the ring with cudaMallocHost as day 53 did.

Reusing a buffer or ring slot without an event wait is a race that a fast card may hide. The correctness pass checks beyond both reuse depths so the race produces a failed byte comparison.

A stray cudaMemcpy or a launch without a stream argument lands on the legacy default stream and can serialize the other streams. Every steady-state call must name its stream; day 51 shows this trace pattern.

A host timestamp taken before the last frame drains reports a rate that is too high. The burst and paced reads follow cudaDeviceSynchronize, the baseline read follows its cudaStreamSynchronize, and generation is timed as host work on an idle device.

At 60 fps, the reference run spends most of each tick idle. Use the paced range to check cadence and the burst range to check overlap at full capacity.

Do not put a cudaDeviceSynchronize between cudaStreamBeginCapture and cudaStreamEndCapture. Capture records launches without running them, so keep only the launches inside and check cudaGetLastError after each.

Sources

CUDA C++ Best Practices Guide: asynchronous transfers, checked 2026-08-29.
CUDA Runtime API: graph management, checked 2026-09-01.
Nsight Systems User Guide, for --trace and nsys stats, checked 2026-08-30.
NVIDIA cuda-samples: cpp/0_Introduction/asyncAPI and cpp/0_Introduction/simpleMultiCopy, checked 2026-09-01.

Day 61 uses compute-sanitizer to find three planted memory bugs. Day 62 then applies racecheck to the missing-event race that this capstone detects with buffer-reuse cases, and day 88 reuses the pipeline structure for serving.