Streams & concurrency

01 · The default serialization

Your GPU has been idling between steps

Every program so far ran on the default stream, and the default stream does one thing religiously: operations execute in launch order, each waiting for the last. So the day-13 ritual — copy in, launch, copy back — executes as a strict relay: the GPU computes nothing while the copy engine fills memory, and PCIe idles while the kernel runs.

// default stream: 3 phases, strictly in order
H2D  ██████████          
KERN            ██████████████
D2H                        ██████████

// two chunks on two streams: copy(k+1) hides under kernel(k)
H2D  ██████████     ██████████
KERN           ██████████████     ██████████████
D2H                      ██████████        ██████████

The second timeline is the opportunity: split the work into chunks and alternate them across two streams, so chunk k+1's transfer overlaps chunk k's compute. If transfer and compute cost similar time, the pipeline approaches 2× throughput — for no change to the math, just the scheduling.

02 · The machinery

Streams, events, and pinned memory

A stream is an ordered queue of device work; work in different streams may overlap. An event is a marker you can record into a stream and wait on — the GPU-side stopwatch. And one prerequisite is easy to miss: async copies need pinned memory. Copying from ordinary pageable RAM forces a staging hop through a locked buffer that silently re-synchronizes the stream — your "async" copy blocks anyway. Allocate staging buffers with cudaMallocHost (page-locked) and the copy engine can truly run behind the CPU's back:

cudaStream_t s0, s1;
cudaStreamCreate(&s0);  cudaStreamCreate(&s1);
float *hbuf;
cudaMallocHost(&hbuf, chunk_bytes);   // pinned: async-capable

for (int k = 0; k < n_chunks; ++k){
    cudaStream_t s = streams[k % 2];
    cudaMemcpyAsync(d_in + off, hbuf, chunk_bytes,
                     cudaMemcpyHostToDevice, s);
    process<<<grid, block, 0, s>>>(d_in + off, d_out + off, m);
    cudaMemcpyAsync(hout + off, d_out + off, chunk_bytes,
                     cudaMemcpyDeviceToHost, s);
}
cudaStreamSynchronize(s0);  cudaStreamSynchronize(s1);

Note the kernel launch grew a third chevron argument: <<<grid, block, smem, stream>>>. Events give you the timing story: record an event before and after, then cudaEventElapsedTime — far more honest than host timers, which measure your CPU's opinion of GPU work.

03 · The honest footnote about copy engines

Consumer cards ship one copy engine

Data-center GPUs have dedicated copy engines that push and pull simultaneously; most GeForce cards — the 4070 included — have a single full-duplex engine, so two host-to-device transfers still queue behind each other. What overlaps reliably is transfer against compute, which is exactly the chunked pattern above. The other overlap that works everywhere: independent kernels on separate streams, when you have two jobs that don't touch the same data.

And the pipeline's shape matters. Overlap helps when transfer time ≈ compute time (day 11's thinking, applied to a timeline instead of a chart). If compute dwarfs transfer, just optimize the kernel; if transfer dwarfs compute, reduce the bytes — day 23's fusion — because no scheduling hides a transfer that's genuinely the critical path.

04 · Where this shows up in real systems

Every training loop is a stream pipeline

Once you see it, the pattern is everywhere: prefetching the next training batch to the GPU while the current one computes; overlapped gradient all-reduce in data-parallel training; inference servers double-buffering requests. The ideas are always the same trio from this chapter — chunk, alternate streams, pin the staging memory. Day 18's Nsight Systems exists precisely to show you these timelines, and once you've built one by hand, its Gantt-style rows will read like the diagram in section 01.

Next: day 18 — reading an Nsight trace: the two profilers, the four numbers to look at first, and the measure-one-change discipline that turns all this knowledge into speed.