Streams: Overlapping Compute and Transfer

A stream is an ordered queue of device work. Operations in the same stream run one after another; operations in different streams may run at the same time. Modern GPUs have separate copy engines, so a transfer on one stream can proceed while a kernel on another stream is computing.

Write a program the obvious way and none of that happens. Everything lands in one stream, and the device spends half its time copying with the arithmetic units idle, then half computing with the copy engines idle.

Every call that does not name a stream goes into the default stream, and the default stream has special serializing behaviour: work in it will not overlap with work in other streams unless you have explicitly opted into a per-thread default. So the standard beginner program, allocate, copy, launch, copy back, is fully sequential by construction.

There is also a trap in cudaMemcpy itself. The plain form is synchronous: it blocks the host until the copy completes. So even creating streams changes nothing if you keep using the blocking copy, because the host never gets far enough ahead to queue the next piece of work.

Three things, and missing any one of them silently gives you no overlap rather than an error.

The pinned memory requirement is the one people miss. An async copy from ordinary malloced memory is async in name only.

float *h_data;
hipHostMalloc(&h_data, bytes);   // pinned; cudaMallocHost is the CUDA spelling
// ... use it ...
hipHostFree(h_data);
            

The payoff is to split the work into chunks and stagger them, so that while chunk 1 is being computed, chunk 2 is being copied in and chunk 0 is being copied out.

const int nStreams = 4;
hipStream_t stream[nStreams];
for (int i = 0; i < nStreams; ++i) hipStreamCreate(&stream[i]);

const int chunk = n / nStreams;
const size_t chunkBytes = chunk * sizeof(float);

for (int i = 0; i < nStreams; ++i) {
    int off = i * chunk;
    hipMemcpyAsync(d_a + off, h_a + off, chunkBytes,
                   hipMemcpyHostToDevice, stream[i]);
    vecAdd<<<chunk / 256, 256, 0, stream[i]>>>(d_a + off, d_b + off, d_c + off, chunk);
    hipMemcpyAsync(h_c + off, d_c + off, chunkBytes,
                   hipMemcpyDeviceToHost, stream[i]);
}

for (int i = 0; i < nStreams; ++i) hipStreamSynchronize(stream[i]);
            

The ceiling on what this buys you is straightforward: if transfer and compute take about the same time, perfect overlap roughly halves the total. If one dominates, you save the smaller of the two and no more.

Events are markers you insert into a stream. They serve two purposes.

hipEvent_t start, stop;
hipEventCreate(&start);
hipEventCreate(&stop);

hipEventRecord(start, stream);
myKernel<<<blocks, threads, 0, stream>>>(args);
hipEventRecord(stop, stream);

hipEventSynchronize(stop);
float ms;
hipEventElapsedTime(&ms, start, stop);
            

For ordering, hipStreamWaitEvent makes one stream wait on an event recorded in another, which expresses a dependency between streams without blocking the host. That is how you build a dependency graph rather than a set of independent pipelines.

Forgetting pinned memory, covered above, is the most common and produces no error.

Calling a synchronizing function in the middle of the loop. A device synchronize, or a plain blocking copy, drains everything and removes the overlap you just built.

Too many streams. Beyond a handful there is nothing left to overlap and you are adding bookkeeping. Four to eight covers most cases.

Assuming it worked. Overlap is invisible from inside the program. Confirm it on a profiler timeline, where overlapping work appears as literally overlapping bars, or not.