CUDA Streams and Concurrency, in depth: default-stream semantics, stream-ordered APIs and libraries, a chunked copy-compute pipeline, host functions and async errors

Stream 0 and how to avoid it

Every API that takes a stream accepts 0 for the default stream, and what that means depends on how the program was built. Under the legacy default stream, stream 0 implicitly synchronises with every other blocking stream on the device: work in it waits for all of them, and they wait for it. A single forgotten launch without a stream argument therefore serialises a carefully overlapped program. Two escapes exist.

Prefer explicit streams everywhere in new code and treat any stream-0 use as a bug unless it is intentional. Streams also have priorities: create them with cudaStreamCreateWithPriority after reading the valid range from cudaDeviceGetStreamPriorityRange, where a numerically lower number means a higher priority. Priority affects which pending blocks the scheduler starts next; it does not preempt blocks already running.

The stream-ordered API surface

The stream is the common currency of the CUDA ecosystem. Any call in this table is asynchronous with respect to the host and ordered within its stream, which is what lets you compose them without host synchronisation.

Two caveats are easy to miss. A copy is only truly asynchronous when the host side is page-locked memory from cudaMallocHost or registered with cudaHostRegister; from pageable memory the driver stages through a bounce buffer and the call may not return until the copy is done. And cudaMalloc and cudaFree are not stream-ordered: cudaFree synchronises the device, so a free in a hot loop quietly removes all your overlap. The stream-ordered allocator and its pools are covered in GPU memory pools.

Writing a stream-correct function

Most stream bugs happen at boundaries between components. A function that launches work should behave like a CUDA API: take the caller's stream, enqueue everything on it, and return without waiting. The rules are short.

1. Accept a stream parameter and use it for every launch, copy, allocation and library call.
2. Set library handles to that stream on every call, because a shared handle may have been pointed at another stream since you last used it.
3. Never call cudaDeviceSynchronize or cudaStreamSynchronize inside; synchronising is the caller's decision.
4. Allocate temporaries with the stream-ordered allocator and free them on the same stream.
5. If you fork work onto internal streams, join them back with events before returning, so the caller's stream is a complete description of your work.

// y = normalize(A * x), entirely ordered on the caller's stream
cudaError_t gemv_normalize(cublasHandle_t h, const float* A, const float* x,
                           float* y, int m, int n, cudaStream_t s)
{
    float* tmp = nullptr;
    cudaError_t err = cudaMallocAsync(&tmp, m * sizeof(float), s);
    if (err != cudaSuccess) return err;

    cublasSetStream(h, s);                         // rule 2: every call
    const float one = 1.0f, zero = 0.0f;
    cublasSgemv(h, CUBLAS_OP_N, m, n, &one, A, m, x, 1, &zero, tmp, 1);

    normalize_kernel<<<(m + 255) / 256, 256, 0, s>>>(tmp, y, m);
    err = cudaGetLastError();                      // launch-time errors only

    cudaFreeAsync(tmp, s);                         // freed after the kernel, in order
    return err;                                    // no sync: caller decides
}

The function never blocks, yet it is safe: the free is ordered after the kernel that reads the buffer, and any later work the caller puts on the same stream sees the finished output. The same discipline makes the function capturable into a CUDA graph without changes, as described in stream capture.

Worked example: a chunked pipeline

The classic use of several streams is splitting a large job into chunks so that copying chunk k + 1 in, computing chunk k and copying chunk k - 1 out happen at once. Suppose 256 MB goes in, a kernel processes it in 8 ms, and 256 MB comes out, over a PCIe link that sustains about 24 GB/s from pinned memory in each direction. Each copy takes about 10.7 ms, so the serial version takes 10.7 + 8 + 10.7 = 29.4 ms.

constexpr int K = 4;                       // chunks, one stream each
cudaStream_t st[K];
for (int k = 0; k < K; ++k)
    cudaStreamCreateWithFlags(&st[k], cudaStreamNonBlocking);

size_t chunk = bytes / K;                  // assume divisible
for (int k = 0; k < K; ++k) {
    size_t off = k * chunk;
    cudaMemcpyAsync(d_in + off, h_in + off, chunk, cudaMemcpyHostToDevice, st[k]);
    process<<<blocks_for(chunk), 256, 0, st[k]>>>(d_in + off, d_out + off, chunk);
    cudaMemcpyAsync(h_out + off, d_out + off, chunk, cudaMemcpyDeviceToHost, st[k]);
}
for (int k = 0; k < K; ++k) cudaStreamSynchronize(st[k]);
// h_in and h_out must come from cudaMallocHost, or nothing overlaps

With K chunks and enough engines, the time approaches the slowest stage plus the other stages divided by K: about 10.7 + (29.4 - 10.7) / 4, roughly 15.4 ms, nearly twice as fast. The count of copy engines matters. Query cudaDeviceProp::asyncEngineCount: with two, uploads and downloads proceed simultaneously as in the figure; with one, both directions share an engine, so the copies alone take at least 21.4 ms, and issue order decides how much of the compute hides inside that. More chunks shrink the fill and drain at the ends but add per-operation overhead, so four to eight chunks is a sensible starting range; measure rather than guess.

Measure with events: record cudaEventRecord(start, st[0]) before the loop, make a final stream wait on every chunk stream, record a stop event there, and read cudaEventElapsedTime. Events used only for ordering should be created with cudaEventDisableTiming, which makes them cheaper.

Host functions in the stream

cudaLaunchHostFunc enqueues a host function that runs after all earlier work in the stream finishes, and later work in that stream waits until it returns. It is the right tool for signalling the host side of a pipeline (for example, marking a pinned staging buffer as free so the loader thread can refill it) without a blocking synchronise. Three rules: the function must not call any CUDA API, it runs on a driver-owned thread so shared state needs proper locking, and it should be short, because the stream stalls while it runs.

struct Slot { std::atomic<bool> free{true}; };

void CUDART_CB release_slot(void* p) {           // no CUDA calls in here
    static_cast<Slot*>(p)->free.store(true, std::memory_order_release);
}

// producer thread
slot.free.store(false);
cudaMemcpyAsync(d_buf, slot_ptr, n, cudaMemcpyHostToDevice, s);
cudaLaunchHostFunc(s, release_slot, &slot);      // fires once the copy has finished

Errors in an asynchronous world

Because calls return before the work runs, errors arrive in two waves. Launch configuration errors (too many threads per block, too much shared memory) are known immediately; check cudaGetLastError right after a launch. Execution errors, such as an out-of-bounds access inside a kernel, are only discovered when the device runs the work, and are reported by whichever later call happens to synchronise or query, possibly in unrelated code. Some of them, an illegal address for example, are sticky: the context is unusable, every subsequent call returns the same error, and the only recovery is to restart the process.

For debugging, set CUDA_LAUNCH_BLOCKING=1 so every launch synchronises and the failing call reports its own error, then turn it off again, because it destroys all overlap. In production, poll with cudaStreamQuery (it returns cudaErrorNotReady while work is pending) or check after each synchronisation point, and treat any sticky error as fatal to the worker.

Streams across several GPUs

Each stream belongs to the device that was current when it was created, and kernels must be launched into a stream of the current device. Events bridge devices: an event recorded on a stream of GPU 0 can be waited on by a stream of GPU 1 with cudaStreamWaitEvent, which is how a multi-GPU program orders a peer copy (cudaMemcpyPeerAsync) after the producer kernel without blocking the host. NCCL works the same way: a collective is just another stream-ordered operation, so overlapping it with compute means giving it its own stream and joining with events. The NCCL collectives article shows the patterns.

Failure modes

To see what actually overlapped, profile with Nsight Systems; the Nsight article explains reading its stream rows.

Trade-offs

Streams buy concurrency at the cost of reasoning: every cross-stream dependency you forget is a race that may pass every test. Use the fewest streams that achieve the overlap you need, usually one for compute, one per copy direction and one for communication. Independent small kernels rarely benefit from extra streams on a large GPU unless each underfills the device. When the launch pattern repeats, record it into a graph to remove per-launch overhead. And keep host synchronisation at the edges of the program: one synchronise at a well-defined point is cheap; many hidden ones are not.

There is also a choice between doing concurrency with streams and doing it inside one kernel. A persistent kernel that loads, computes and stores its own tiles with asynchronous copies can overlap memory and arithmetic at a much finer grain than chunked streams, with no host involvement at all, but it is harder to write and ties the overlap to one kernel's code. Streams work at the granularity of whole operations and compose across libraries you did not write. Use streams to overlap independent operations and transfers; reach for in-kernel pipelining only when profiling shows the remaining stall is inside a single kernel. Either way, verify the result in a profile: overlap that is only assumed is often missing.

What to do next

1. Search your code for launches and copies without a stream argument and decide, for each, whether stream 0 is intended.
2. Make every host buffer used with async copies pinned, and verify it in a profile.
3. Refactor one GPU helper to the library rules above: take a stream, never synchronise, allocate stream-ordered.
4. Build the four-chunk pipeline for one real transfer, read asyncEngineCount, and compare the measured time against the formula.
5. Check cudaGetLastError after launches and route sticky errors to a worker restart.
6. Move on to stream capture once the stream structure is stable.