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

CUDA streams and overlap

A learner on NVIDIA's forum launches three kernels and a copy, expects them to run together, and watches them refuse. The answer they get is blunt: "kernel1, kernel2, kernel3, and cudaMemcpy will not run in parallel, but sequentially." (https://forums.developer.nvidia.com/t/cuda-kernel-synchronization-issue-host-cudamemcpy-stuck-high-gpu-utilization/322076 , checked 2026-08-29).

The intuition behind the disappointment is common: the GPU is a parallel machine, so surely two independent kernels run at once. By default they do not, and when you fix the default they still might not.

By the end of this page you will know the two separate things overlap requires. You will have a timeline showing two kernels running together and know why the most obvious version of the experiment, two big kernels in two streams, shows nothing at all.

What a stream is, and what overlap actually needs

A stream is an ordered queue of GPU work. Within one stream, operation N+1 starts after operation N finishes. Between two streams, CUDA promises nothing, and that absence of a promise is the whole feature: it is permission to overlap, not a command to.

Every launch so far in this course went to the same queue. Leave the fourth launch parameter empty and work lands in the legacy default stream, which the runtime API documents as "an implicit stream which synchronizes with all other streams in the same CUcontext except for non-blocking streams" (https://docs.nvidia.com/cuda/cuda-runtime-api/stream-sync-behavior.html , checked 2026-09-01). One queue, everything ordered, which is why 50 days of lessons never needed this page.

Permission is the first requirement. The second is spare hardware. A kernel launch hands the GPU a grid of thread blocks, and the scheduler places blocks on SMs until it runs out of blocks or room.

A second kernel's blocks can start only where the first kernel leaves room. If the first grid fills every SM, the second waits. A small first grid may leave enough SMs for work from another stream.

So overlap needs both: independent work in separate streams (permission) and a card that is not already full (room). The widget below is the model to keep. Its one-stream preset is every program you have written so far; two-streams-pinned is the shape this module builds toward, copies and kernels running at the same time.

Streams do not create parallelism, they permit it

Creating a stream does not make work faster by itself. A stream records order and changes only what the GPU may run at the same time. If the work was going to saturate the card anyway, streams rearrange the queue and the wall clock does not move.

The defaults can also serialize work before the hardware limit does. cudaStreamCreate gives you a blocking stream: your kernel runs in it, but anything in the legacy default stream still serialises against it. The mechanism is explicit in the documentation: "When an action is taken in the legacy stream such as a kernel launch or cudaStreamWaitEvent(), the legacy stream first waits on all blocking streams, the action is queued in the legacy stream, and then all blocking streams wait on the legacy stream" (same page as above, checked 2026-09-01).

One forgotten launch with an empty fourth argument, or one library call that uses the default stream, makes the program sequential again. Streams created with the cudaStreamNonBlocking flag do not synchronize with the legacy stream. nvcc --default-stream per-thread changes the default itself; this lesson keeps the legacy default and shows the behavior directly.

One binary, seven arrangements

Full program in code/day51-streams/streams.cu. Four rules hold it together.

Only the launch arrangement moves. Two kernels: spinLcg, 8 blocks of dependent integer arithmetic that hold at most 8 SMs for milliseconds while barely touching memory, and sweepAdd, a 320-block bandwidth sweep that fills the card. The program times each alone, then both in the legacy stream, then in two non-blocking streams, then the blocking-stream trap, then two grid-filling sweeps against each other. It uses the same kernels every time.

Correctness before timing. All three kernels (there is a sweepScale twin so the two lanes of part 4 have distinct names) run once against exact CPU references before anything is timed. The inputs are small whole numbers, so the float arithmetic is exact and any mismatch is a launch bug.

Events time it, the timeline proves it. Each arrangement is a closure ending in cudaDeviceSynchronize, so the async work is drained inside the closure, before the stop event is recorded. That is where the synchronise lives, and it is why the event pair brackets real work instead of launch overhead, the mistake day 9 priced.

Every arrangement is an NVTX range, using the same name as the results table, per day 41's conventions, so nsys stats --report nvtx_sum and the program's own table line up by name.

Here are the two kernels:

__global__ void spinLcg(unsigned int* out, size_t n, int iters) {
    const size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
    if (i < n) {
        unsigned int x = static_cast<unsigned int>(i);
        for (int k = 0; k < iters; ++k) {
            x = 1664525u * x + 1013904223u;
        }
        out[i] = x;
    }
}

// Grid-stride sweep: out[i] = in[i] + 1.0f. Bandwidth-bound, and with
// kFullBlocks blocks it fills every SM. Consecutive threads read and write
// consecutive floats, so each warp moves 128-byte contiguous lines.
__global__ void sweepAdd(const float* in, float* out, size_t n) {
    const size_t step = gridDim.x * static_cast<size_t>(blockDim.x);
    for (size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
         i < n; i += step) {
        out[i] = in[i] + 1.0f;
    }
}

And the arrangement the page is named for. The stream is the fourth chevron argument; the third, dynamic shared memory from day 13, is written out only to reach it:

        auto part2TwoStreams = [&] {
            nvtxRangePushA("part2-two-streams");
            spinLcg<<<kSpinBlocks, kThreadsPerBlock, 0, streamA>>>(
                d_spin, kSpinThreads, kSpinIters);
            for (int s = 0; s < kSweepsPerLane; ++s) {
                sweepAdd<<<kFullBlocks, kThreadsPerBlock, 0, streamB>>>(
                    d_in, d_a, kElems);
            }
            CUDA_CHECK(cudaDeviceSynchronize());
            nvtxRangePop();
        };

Part 3 differs by two tokens. streamC came from plain cudaStreamCreate, and the sweeps go to the legacy default stream:

        auto part3LegacyTrap = [&] {
            nvtxRangePushA("part3-legacy-trap");
            spinLcg<<<kSpinBlocks, kThreadsPerBlock, 0, streamC>>>(
                d_spin, kSpinThreads, kSpinIters);
            for (int s = 0; s < kSweepsPerLane; ++s) {
                sweepAdd<<<kFullBlocks, kThreadsPerBlock>>>(d_in, d_a, kElems);
            }
            CUDA_CHECK(cudaDeviceSynchronize());
            nvtxRangePop();
        };

Note. The timer in this program can only ever corroborate. A wall clock cannot distinguish "the kernels overlapped" from "the second kernel was faster than I thought". The proof of overlap is the Nsight Systems timeline, where each stream is its own row, or the cuda_gpu_trace export, where overlap is two rows with different stream ids and intersecting intervals. Text first, per day 41.

Results

Re-verified 2026-09-02 on the same Tesla T4 with driver 580.173.02 and CUDA 13.0 (V13.0.88). Correctness and all three qualitative outcomes reproduced: the useful overlap still won, the legacy-stream trap still did not, and two full grids still cost their solo sum.

The CUDA 13 transcript and fresh Nsight Systems report are listed in front matter; the original CUDA 12.6 transcript and published report remain as provenance.

The shipped report and CSV exports live at content/profiles/cuda-streams/: day51-streams.nsys-rep with its nvtx_sum, cuda_gpu_kern_sum and cuda_gpu_trace exports beside it; the exercise is completable from them alone. The card reports 3 async copy engines and holds 160 of part 4's 320 blocks resident.

spin-solo                  2.156 ms  (min    2.154, max    2.159)
sweep-solo                 2.545 ms  (min    2.535, max    2.555)
scale-solo                 2.573 ms  (min    2.564, max    2.596)
part1-legacy               4.154 ms  (min    4.141, max    4.172)
part2-two-streams          2.719 ms  (min    2.712, max    2.724)
part3-legacy-trap          4.113 ms  (min    4.094, max    4.155)
part4-two-full-grids       5.085 ms  (min    5.080, max    5.088)

part1 / part2 (overlap saving)       1.53x
part2 / spin-solo (hiding quality)   1.26x
part3 / part1 (trap vs sequential)   0.99x
part4 / (sweep + scale solo)         0.99x

What I expect, written down before the run

  1. Part 2 lands within 10 percent of spin-solo, and part 1 costs at least 1.3 times part 2. The four sweepAdd bars should sit inside the one spinLcg bar on a different stream row of the timeline, and in cuda_gpu_trace their intervals should intersect spinLcg's. That row is the artifact this page exists for.
  2. Part 3 lands within 10 percent of part 1. A second stream, zero benefit: cudaStreamCreate makes a blocking stream, and the legacy launches serialise against it exactly as the quoted rule says.
  3. Part 4's printed ratio against the sum of its two solo lanes lands at 0.9 or above. Real overlap would drag it toward 0.5. Both sweeps need all 40 SMs and the same memory bandwidth, and neither can lend what it is using. Whether the timeline shows strict turn-taking or wave-level interleaving of the two grids is the one detail I refuse to predict; the wall-clock claim stands either way.
  4. In cuda_gpu_trace, the five kernels of one part2-two-streams pass together run for longer than the wall interval they span. Sum the five durations, subtract the earliest start from the latest end; the sum comes out larger, which is impossible without concurrency. The pass is findable in the trace with no GUI open: one spinLcg and four sweepAdd rows on two stream ids, intervals intersecting.

Prediction 3 tests the model most directly. If two grid-filling kernels do overlap productively on this card, then "spare SMs" is the wrong model for what overlap needs, and the module's framing gets rewritten against whatever the timeline shows.

What the run said

  1. Half held; the hiding bound missed more clearly. Part 1 costs 1.53x part 2, well past the 1.3 bar. But part 2 landed at 1.11x spin-solo, one point outside the 10 percent bound in CUDA 12.6 and at 1.26x in CUDA 13.0: the four sweeps do not hide for free under the spin, and now cost 0.563 ms of visible time. The bound was wrong, not the mechanism; the trace shows the overlap.
  2. Held. Part 3 came in at 0.99x part 1, inside the 10 percent band, and on the trap's side of it: a second stream bought nothing next to legacy-stream launches.
  3. Held. Part 4's ratio printed 0.99: two full grids in two streams cost the sum of their solos, almost to the microsecond.
  4. Held in the published CUDA 12.6 trace, and it is the proof. In cuda_gpu_trace, one part2-two-streams pass is one spinLcg and four sweepAdd rows on stream ids 13 and 14: the five durations sum to 4.742 ms while the interval they span is 2.773 ms. The kernels ran 1.71x more kernel-time than wall-time, which no sequential schedule can do.

A stream row that starts while another is busy is the only proof of overlap. Whatever the numbers say, do not accept a faster total as evidence by itself.

Run it yourself

Use a CUDA GPU with Nsight Systems installed.

nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo -o streams streams.cu
nsys profile -t cuda,nvtx -o day51-streams ./streams
nsys stats --report nvtx_sum --report cuda_gpu_kern_sum \
    --report cuda_gpu_trace --format csv \
    --output day51-streams day51-streams.nsys-rep

No root and no ncu: overlap is a timeline question, and nsys traces as a normal user. There is no Compiler Explorer embed because no profiler runs there and the proof is the timeline, not the program's own output. If you have no GPU, read /setup/learn-cuda-without-a-gpu and work from the shipped reports at content/profiles/cuda-streams/.

Exercise

Kill the overlap without touching a stream. Change kSpinBlocks from 8 to 320, predict what happens to the part 1 to part 2 ratio before you run, then run and check your prediction on the timeline.

Time: 20 to 35 minutes. Submit: your predicted and measured part1 / part2 ratios at 320 spin blocks, and one sentence naming the resource that changed.

Check: compare the ratio, not the absolute time, so results remain useful across GPUs. At 8 spin blocks the program's part1 / part2 line should show a real saving; at 320 it should collapse to within 10 percent of 1.0. If it does not collapse, either your card has enough SMs that 320 blocks of spin still leave room for the sweeps (check the resident-block arithmetic the program prints at startup), or the kernels found another shared resource, and the timeline will show which.

Hint 1

Nothing about the streams changed, so whatever permission they granted is still granted. What did the small kernel stop leaving behind when it grew?

Hint 2

The program prints the arithmetic at startup: resident threads per SM divided by the block size, times the SM count, is how many blocks fill the card. Compare 320 spin blocks against that number, then look at where the sweep bars start on the kernel row.

Solution

At 320 blocks, spinLcg fills every SM on the reference card (160 resident blocks fill it; the rest queue behind them), so the sweeps in stream B have permission to run but no free SMs. Their bars should start after the spin bar, and part 2's time should climb to roughly part 1's, so the ratio falls toward 1.0. The streams are unchanged and working; the room is gone.

Overlap needs both stream permission and free hardware. When streams "stop working", check whether the first kernel leaves any SMs free.

Pitfalls

You created streams and the timeline is still sequential. Either a launch's fourth argument is empty, so it went to the legacy default stream, or your streams came from plain cudaStreamCreate and something else is using the legacy stream. Both serialise, per the legacy-stream rule quoted above.

Create streams with cudaStreamNonBlocking, or build with nvcc --default-stream per-thread. This page's part 3 measures the trap.

Two big kernels in two streams do not overlap. Not a bug. Each grid fills the card, so the scheduler has no free SM to place the other kernel's blocks on.

Size one kernel down and the overlap appears, which is this page's part 2 against part 4. Day 45 is the arithmetic of how much of an SM a block occupies.

Your event timer says the streamed version is hundreds of times faster. The launches are asynchronous, so an event pair with no synchronise in front of the stop records queueing, not work. Synchronise before the stop event and say so in a comment; this program does it inside each timed closure.

The copies in your pipeline still refuse to overlap. Kernel overlap is today's half of the story. A copy overlaps compute only from pinned host memory with cudaMemcpyAsync in a non-default stream: "the overlap once again requires pinned host memory, and, in addition, the data transfer and kernel must use different, non-default streams" (https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html , checked 2026-09-01). Day 53 measures it.

You profiled with Nsight Compute and the overlap vanished. Its profiling guide says plainly that "Nsight Compute serializes kernel launches, unless a dedicated replay mode is used" (https://docs.nvidia.com/nsight-compute/ProfilingGuide/index.html , checked 2026-09-01). Concurrency questions go to nsys; per-kernel questions go to ncu. Using the wrong one does not just miss the answer, it changes it.

Go deeper

Next

Day 52 takes the one tool this page did not need, the event, and uses it for what it is really for: making stream B wait for one specific point in stream A instead of for everything, which turns two independent queues into a fork and a join. Day 53 then adds the copy engines, where day 9's expensive round trip starts earning its keep, and day 48's launch-cost arithmetic decides how fine to slice the work.