Streams: What They Are and How to Create Them

By the end of this page you can create and destroy streams, explain why a single queue can never overlap anything, and check whether your GPU is even capable of overlap.

What a stream is

A stream is an in-order queue of device operations. Everything you put into one stream — copies, kernel launches, event records — executes in the order you submitted it, and the next operation doesn’t start until the previous one finishes. From the CUDA C Programming Guide (12.8): “A stream is a sequence of commands (possibly issued by different host threads) that execute in order. Different streams, on the other hand, may execute their commands out of order with respect to one another or concurrently.”

That last sentence is the whole point. One stream = one queue = strict serialization. Two or more streams = independent queues that the driver can run concurrently when the hardware has idle resources. Copy engines and SMs are different resources, so a copy in stream A and a kernel in stream B can genuinely run at the same time. That is the overlap you’re after, and it is impossible with a single stream.

The default stream (the NULL stream) and why it bites you

Kernel launches and copies that don’t name a stream go to the default stream (stream 0, the NULL stream). Here’s the trap: with default compilation flags (legacy behavior), the default stream is special — any operation on it implicitly synchronizes with every blocking stream. From the guide: “Two operations from different streams cannot run concurrently if any CUDA operation on the NULL stream is submitted in-between them, unless the streams are non-blocking streams (created with the cudaStreamNonBlocking flag).”

In practice: a stray cudaMemcpy (no stream) between your streamed operations can silently serialize everything around it, killing your overlap. Two habits fix this:

1. Always create your streams with cudaStreamNonBlocking, so they never participate in default-stream implicit synchronization.
2. Never use the default stream at all inside code you want to overlap. (If you compile with --default-stream per-thread, the default stream becomes an ordinary per-thread stream — a compiler flag worth knowing, but this guide just uses explicit non-blocking streams.)

The one legitimate use of the default stream: a quick synchronous cudaMemcpy for correctness-critical setup, where implicit sync is exactly what you want.

Creating and destroying streams

Signatures verified against cuda_runtime_api.h (CUDA 13.2):

cudaError_t cudaStreamCreate(cudaStream_t *pStream);
cudaError_t cudaStreamCreateWithFlags(cudaStream_t *pStream, unsigned int flags);
cudaError_t cudaStreamDestroy(cudaStream_t stream);

cudaStreamCreate is equivalent to cudaStreamCreateWithFlags(&s, cudaStreamDefault) (0x00). The flag that matters is cudaStreamNonBlocking (0x01) — from the docs: “Specifies that work running in the created stream may run concurrently with work in stream 0 (the NULL stream), and that the created stream should perform no implicit synchronization with stream 0.” Use it for every stream you create in this guide.

cudaStreamDestroy is non-blocking: it returns immediately, and the stream’s resources are released automatically once the device finishes the work still queued in it. You do not need to drain a stream before destroying it (though you should before freeing buffers it might still be using).

Does your GPU even support overlap?

Two device properties decide it (fields of cudaDeviceProp, verified in driver_types.h):

• asyncEngineCount — “Number of asynchronous engines”. Greater than 0 means the device can copy to/from the GPU concurrently with kernel execution. Equal to 2 means host-to-device and device-to-host copies can also overlap each other. This is the property that gates everything in this guide.
• concurrentKernels — 1 means multiple kernels can execute concurrently.

If asyncEngineCount is 0 (some integrated and very old GPUs), streams still work — they just serialize, and page 04’s benchmark will show no speedup. That’s expected, not a bug.

The program

stream_basics.cu — prints your GPU’s overlap capabilities and creates/destroys two non-blocking streams:

// stream_basics.cu — device properties and stream creation (CUDA Toolkit 13.2)
// Compile:  nvcc -arch=sm_80 -O2 stream_basics.cu -o stream_basics
#include <cstdio>
#include <cstdlib>

#define CUDA_CHECK(call)                                                        \
  do {                                                                          \
    cudaError_t err = (call);                                                   \
    if (err != cudaSuccess) {                                                   \
      fprintf(stderr, "CUDA error at %s:%d: %s\n", __FILE__, __LINE__,          \
              cudaGetErrorString(err));                                         \
      exit(1);                                                                  \
    }                                                                           \
  } while (0)

int main() {
  int devCount = 0;
  CUDA_CHECK(cudaGetDeviceCount(&devCount));
  if (devCount == 0) {
    fprintf(stderr, "no CUDA-capable device is detected\n");
    return 1;
  }

  cudaDeviceProp prop;
  CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));
  printf("GPU:                %s\n", prop.name);
  printf("asyncEngineCount:   %d  (DMA copy engines; 0 = no copy/kernel overlap)\n",
         prop.asyncEngineCount);
  printf("concurrentKernels:  %d  (1 = multiple kernels can run at once)\n",
         prop.concurrentKernels);

  // cudaStreamNonBlocking: this stream does NOT implicitly synchronize
  // with the NULL (default) stream — see the guide page on streams.
  cudaStream_t s1, s2;
  CUDA_CHECK(cudaStreamCreateWithFlags(&s1, cudaStreamNonBlocking));
  CUDA_CHECK(cudaStreamCreateWithFlags(&s2, cudaStreamNonBlocking));
  printf("created non-blocking streams: %p, %p\n", (void *)s1, (void *)s2);

  CUDA_CHECK(cudaStreamDestroy(s1));
  CUDA_CHECK(cudaStreamDestroy(s2));
  printf("streams destroyed\n");
  return 0;
}

The CUDA_CHECK macro is worth keeping: async APIs report failures through error codes, and errors from previous async launches can surface on later calls. Checking every call is the difference between a silent wrong answer and a one-line diagnosis.

Compile (real, verified — nvcc 13.2.86):

nvcc -arch=sm_80 -O2 stream_basics.cu -o stream_basics

Pick -arch for your GPU: sm_80 = A100 / compute 8.0, sm_86 = GeForce RTX 30-series, sm_89 = RTX 40-series, sm_90 = H100. When in doubt, nvidia-smi shows your compute capability.

What running it on a machine with the toolkit but no NVIDIA driver produces (real captured output — this is exactly what this guide’s build box prints):

CUDA error at stream_basics.cu:18: CUDA driver version is insufficient for CUDA runtime version

That’s cudaErrorInsufficientDriver, and it’s the standard “driver missing entirely” failure. On a machine with a driver but no GPU you instead get (verified error strings):

no CUDA-capable device is detected

On a real GPU box the success path prints, e.g.:

GPU:                NVIDIA GeForce RTX 3090
asyncEngineCount:   2  (DMA copy engines; 0 = no copy/kernel overlap)
concurrentKernels:  1  (1 = multiple kernels can run at once)
created non-blocking streams: 0x55f0..., 0x55f1...
streams destroyed

(Expected output — not executed here, no GPU on the build host.)

Version traps (both real, both hit while building this guide)

On a 2025-or-newer distro (this site’s build host runs Ubuntu 26.04), CUDA 12.9 fails twice before you get to write any code:

error: #error -- unsupported GNU version! gcc versions later than 14 are not supported!

Then, even with gcc-14 installed (-ccbin gcc-14), the new glibc headers break the build:

/usr/include/x86_64-linux-gnu/bits/mathcalls.h(83): error: exception specification is
incompatible with that of previous function "cospi" (declared at line 2601 of .../crt/math_functions.h)

Both are real captures. The clean fix on modern distros: install CUDA 13.x (cuda-nvcc-13-2 on Ubuntu), which supports current gcc and glibc. Don’t fight it with -allow-unsupported-compiler.

The silent stream default

cudaMemcpyAsync(..., stream) has stream = 0 as a default argument — omitting it compiles with only a warning (verified: “warning: variable ‘d’ is used before its value is set”, and the call binds to the default). You just put a copy on the default stream, which implicitly synchronizes with your non-blocking streams and can serialize your whole pipeline. Always pass the stream explicitly.