Streams, Events & Asynchronous Execution

A stream is a queue of operations (kernels, copies) executed in order; operations in different streams may run concurrently.

17.1 Streams

A stream is a queue of operations (kernels, copies) executed in order; operations in different streams may run concurrently.

The default (legacy) stream synchronizes with all other blocking streams; compile with --default-stream per-thread to give each host thread its own default stream.

cudaStream_t s1, s2;
cudaStreamCreate(&s1);
cudaStreamCreateWithFlags(&s2, cudaStreamNonBlocking);   // doesn't sync with the legacy default stream

kernelA<<<grid, block, 0, s1>>>(...);
kernelB<<<grid, block, 0, s2>>>(...);                   // may overlap with kernelA
cudaMemcpyAsync(d, h, bytes, cudaMemcpyHostToDevice, s1);
cudaStreamSynchronize(s1);
cudaStreamQuery(s2);                                     // non-blocking status check
cudaStreamDestroy(s1); cudaStreamDestroy(s2);

17.2 Overlapping transfers and computation

Requirements: pinned host memory (§18), GPUs with copy engines (most have 2+ for bidirectional overlap). Pipeline: H2D copy of chunk k+1 ‖ kernel on chunk k ‖ D2H copy of chunk k−1.

const int nStreams = 4;
int chunk = n / nStreams;
cudaStream_t streams[nStreams];
for (int s = 0; s < nStreams; ++s) cudaStreamCreate(&streams[s]);

for (int s = 0; s < nStreams; ++s) {
    int off = s * chunk;
    cudaMemcpyAsync(d_in + off, h_in + off, chunk * sizeof(float), cudaMemcpyHostToDevice, streams[s]);
    process<<<(chunk + 255) / 256, 256, 0, streams[s]>>>(d_in + off, d_out + off, chunk);
    cudaMemcpyAsync(h_out + off, d_out + off, chunk * sizeof(float), cudaMemcpyDeviceToHost, streams[s]);
}
cudaDeviceSynchronize();

17.3 Events — timing & cross-stream dependencies

cudaEvent_t start, stop;
cudaEventCreate(&start); cudaEventCreate(&stop);
cudaEventRecord(start, stream);
kernel<<<grid, block, 0, stream>>>(...);
cudaEventRecord(stop, stream);
cudaEventSynchronize(stop);
float ms; cudaEventElapsedTime(&ms, start, stop);
printf("Kernel time: %.3f ms\n", ms);

// Make stream s2 wait for work recorded in s1
cudaEvent_t ev; cudaEventCreateWithFlags(&ev, cudaEventDisableTiming);
cudaEventRecord(ev, s1);
cudaStreamWaitEvent(s2, ev, 0);

17.4 Host callbacks & priorities

cudaLaunchHostFunc(stream, [](void* data) { printf("chunk done\n"); }, nullptr);   // runs on a host thread
int lo, hi; cudaDeviceGetStreamPriorityRange(&lo, &hi);
cudaStreamCreateWithPriority(&urgent, cudaStreamNonBlocking, hi);

17.5 Concurrent kernels

Multiple small kernels from different streams can run simultaneously if resources allow (Hyper-Q: 32 hardware queues). Visualize overlap with Nsight Systems.