CUDA C++: Streams, Async Copies and Overlap

A stream is a queue of GPU work. Copies and kernels in one stream run in order; work in different streams can run at the same time, so a copy can move the next chunk of data while a kernel computes on the last one. This lesson shows the calls returning early, the memory that quietly makes them synchronous, a pipeline that overlaps copies with compute, and the two ways streams wait for each other when you didn’t ask them to. Every program ran on a Tesla T4 with CUDA 13.0 (driver 580.82.07, nvcc V13.0.88).

Asynchronous calls return before the work is done

Here a 64 MB copy to the GPU, a kernel and a copy back are queued on one stream, and the CPU checks the clock before and after waiting for them.

// async_return.cu - a copy, a kernel and a copy back on one stream return before the GPU is done
#include <chrono>
#include <cstdio>
#include <cuda_runtime.h>

__global__ void work(float* x, int n, int iters) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) { float v = x[i]; for (int k = 0; k < iters; k++) v = v * 0.5f + 1.0f; x[i] = v; }
}

double ms_since(std::chrono::steady_clock::time_point t0) {
  return std::chrono::duration<double, std::milli>(std::chrono::steady_clock::now() - t0).count();
}

int main() {
  const int n = 1 << 24;  // 16M floats, 64 MB
  float *h, *d;
  cudaMallocHost(&h, n * 4);  // page-locked host memory (more on this below)
  cudaMalloc(&d, n * 4);
  for (int i = 0; i < n; i++) h[i] = i % 100;
  cudaStream_t s;
  cudaStreamCreate(&s);
  work<<<1, 1, 0, s>>>(d, 1, 1);  // warm-up: loads the kernel (CUDA loads kernels lazily, see below)
  cudaStreamSynchronize(s);

  auto t0 = std::chrono::steady_clock::now();
  cudaMemcpyAsync(d, h, n * 4, cudaMemcpyHostToDevice, s);
  work<<<n / 256, 256, 0, s>>>(d, n, 1000);
  cudaMemcpyAsync(h, d, n * 4, cudaMemcpyDeviceToHost, s);
  printf("all three calls returned after %6.2f ms\n", ms_since(t0));
  printf("cudaStreamQuery right away:    %s\n", cudaGetErrorName(cudaStreamQuery(s)));
  cudaStreamSynchronize(s);
  printf("cudaStreamSynchronize done at  %6.2f ms\n", ms_since(t0));
  printf("cudaStreamQuery after sync:    %s\n", cudaGetErrorName(cudaStreamQuery(s)));

  int wrong = 0;  // CPU reference: v = v * 0.5 + 1 converges to exactly 2.0
  for (int i = 0; i < n; i++) wrong += h[i] != 2.0f;
  printf("elements that differ from the CPU answer: %d of %d\n", wrong, n);
  cudaStreamDestroy(s); cudaFreeHost(h); cudaFree(d);
}

$ nvcc -o async_return async_return.cu && ./async_return
all three calls returned after   0.02 ms
cudaStreamQuery right away:    cudaErrorNotReady
cudaStreamSynchronize done at   17.85 ms
cudaStreamQuery after sync:    cudaSuccess
elements that differ from the CPU answer: 0 of 16777216

All three calls returned in 0.02 ms, but the work took 17.9 ms (three runs: 0.02 to 0.03 ms, and 14.8 to 17.9 ms). cudaMemcpyAsync and a launch with a stream argument (<<<grid, block, 0, s>>>) only put work in the stream’s queue. The GPU runs that queue in order: the kernel can’t start until the copy in has finished, and the copy back waits for the kernel.

cudaStreamQuery asks whether the queue is empty without waiting. cudaErrorNotReady isn’t a failure, it means “still running”. cudaStreamSynchronize blocks the CPU until the queue is empty, and only after it is the result safe to read: every element matches the CPU’s answer of 2.0, which v * 0.5 + 1 reaches exactly from any starting value here.

Pageable memory makes an async copy synchronous

The example above used cudaMallocHost. This one copies the same 256 MB twice, from ordinary malloc memory and from cudaMallocHost memory, and times both the call and the copy.

// pageable_trap.cu - cudaMemcpyAsync from malloc'd memory blocks the CPU; from pinned memory it doesn't
#include <chrono>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <cuda_runtime.h>

void copy_from(const char* label, float* h, float* d, size_t bytes, cudaStream_t s) {
  cudaEvent_t a, b; cudaEventCreate(&a); cudaEventCreate(&b);
  auto t0 = std::chrono::steady_clock::now();
  cudaEventRecord(a, s);
  cudaMemcpyAsync(d, h, bytes, cudaMemcpyHostToDevice, s);
  cudaEventRecord(b, s);
  double ret = std::chrono::duration<double, std::milli>(std::chrono::steady_clock::now() - t0).count();
  cudaStreamSynchronize(s);
  float gpu; cudaEventElapsedTime(&gpu, a, b);
  printf("%-8s call returned after %7.2f ms, copy took %6.2f ms (%5.2f GB/s)\n",
         label, ret, gpu, bytes / gpu / 1e6);
}

int main() {
  const size_t n = 1 << 26, bytes = n * 4;  // 256 MB
  float *pageable = (float*)malloc(bytes), *pinned, *d, *back = (float*)malloc(bytes);
  cudaMallocHost(&pinned, bytes);
  cudaMalloc(&d, bytes);
  for (size_t i = 0; i < n; i++) pageable[i] = pinned[i] = i % 1000;
  cudaStream_t s; cudaStreamCreate(&s);
  copy_from("warm-up", pinned, d, bytes, s);
  for (int run = 0; run < 2; run++) {
    copy_from("malloc", pageable, d, bytes, s);
    copy_from("pinned", pinned, d, bytes, s);
  }
  cudaMemcpy(back, d, bytes, cudaMemcpyDeviceToHost);
  printf("device copy matches the host data: %s\n", memcmp(back, pageable, bytes) == 0 ? "yes" : "no");
}

$ nvcc -o pageable_trap pageable_trap.cu && ./pageable_trap
warm-up  call returned after    0.06 ms, copy took  21.71 ms (12.37 GB/s)
malloc   call returned after   58.21 ms, copy took  58.28 ms ( 4.61 GB/s)
pinned   call returned after    0.01 ms, copy took  21.70 ms (12.37 GB/s)
malloc   call returned after   58.42 ms, copy took  58.46 ms ( 4.59 GB/s)
pinned   call returned after    0.02 ms, copy took  21.70 ms (12.37 GB/s)
device copy matches the host data: yes

From malloc memory, cudaMemcpyAsync held the CPU for the whole copy, 57 to 61 ms in four runs, and moved data at 4.4 to 4.7 GB/s. From pinned memory the call returned in 0.01 or 0.02 ms and the copy ran at 12.35 to 12.37 GB/s.

The GPU’s copy engines read host memory directly, which needs page-locked (pinned) memory: pages the operating system promises not to move or swap out. malloc memory isn’t pinned, so the driver copies it into a pinned staging buffer first, using the CPU, and the CUDA runtime docs say the call returns only once the data has reached that staging buffer. For a large copy that’s almost all of it. The extra pass through the staging buffer is also why it’s slower.

The fix is to allocate the buffers you copy from with cudaMallocHost (or cudaHostAlloc, or pin existing memory with cudaHostRegister). Pinned memory is taken away from the operating system, so pin the buffers you stream through, not everything. The Python version of this comparison is in CUDA Python: memory transfers.

Overlapping copies and compute across streams

With pinned memory, the copies and kernels of different streams can run at the same time. This program splits 64 MB into 8 chunks; each chunk gets a copy in, a kernel and a copy out, and chunk c goes to stream c % nstreams.

// pipeline.cu - copy in, compute, copy out in 8 chunks, spread over 1, 2 or 4 streams
#include <chrono>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <cuda_profiler_api.h>
#include <cuda_runtime.h>

__global__ void work(float* x, int n, int iters) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) {
    float v = x[i];
    for (int k = 0; k < iters; k++) v = v * 0.5f + 1.0f;  // converges to 2.0
    x[i] = v + x[i];                                       // so the answer is 2 + input
  }
}

const int N = 1 << 24, CHUNKS = 8, C = N / CHUNKS;

double run(int nstreams, float* in, float* out, float* d, cudaStream_t* s) {
  memset(out, 0, N * 4);
  auto t0 = std::chrono::steady_clock::now();
  for (int c = 0; c < CHUNKS; c++) {
    cudaStream_t st = s[c % nstreams];  // chunk c goes to stream c mod nstreams
    size_t off = (size_t)c * C;
    cudaMemcpyAsync(d + off, in + off, C * 4, cudaMemcpyHostToDevice, st);
    work<<<C / 256, 256, 0, st>>>(d + off, C, 1000);
    cudaMemcpyAsync(out + off, d + off, C * 4, cudaMemcpyDeviceToHost, st);
  }
  cudaDeviceSynchronize();
  return std::chrono::duration<double, std::milli>(std::chrono::steady_clock::now() - t0).count();
}

int main(int argc, char** argv) {
  int only = argc > 1 ? atoi(argv[1]) : 0;  // ./pipeline 4: one warm-up, then one profiled run
  cudaDeviceProp p; cudaGetDeviceProperties(&p, 0);
  if (!only) printf("%s, asyncEngineCount = %d\n", p.name, p.asyncEngineCount);
  float *in, *out, *d;
  cudaMallocHost(&in, N * 4); cudaMallocHost(&out, N * 4); cudaMalloc(&d, N * 4);
  for (int i = 0; i < N; i++) in[i] = i % 100;
  cudaStream_t s[4];
  for (int i = 0; i < 4; i++) cudaStreamCreate(&s[i]);
  if (only) {
    run(only, in, out, d, s);
    cudaProfilerStart(); run(only, in, out, d, s); cudaProfilerStop();  // nsys records only this
    return 0;
  }
  run(4, in, out, d, s);  // warm-up
  for (int ns : {1, 2, 4}) {
    double total = 0;
    for (int rep = 0; rep < 5; rep++) total += run(ns, in, out, d, s);
    int wrong = 0;
    for (int i = 0; i < N; i++) wrong += out[i] != 2.0f + i % 100;
    printf("%d stream(s): %6.2f ms (mean of 5), wrong elements: %d\n", ns, total / 5, wrong);
  }
}

$ nvcc -o pipeline pipeline.cu && ./pipeline
Tesla T4, asyncEngineCount = 3
1 stream(s):  19.05 ms (mean of 5), wrong elements: 0
2 stream(s):  10.63 ms (mean of 5), wrong elements: 0
4 stream(s):   9.19 ms (mean of 5), wrong elements: 0

The same 24 operations took about half as long on two or four streams. Across four runs, one stream took 16.7 to 20.9 ms, two streams 9.6 to 11.7 ms and four streams 7.9 to 11.3 ms; every run had 0 wrong elements. On one stream each chunk waits for the previous chunk’s copy out. With more streams, chunk 1’s copy in can run while chunk 0’s kernel computes.

asyncEngineCount is the number of asynchronous copy engines. With 2 or more, the GPU can copy in, copy out and run a kernel all at once. Two streams already capture most of the gain here, because the kernels are the longest stage: once the copies are hidden behind them, more streams have nothing left to hide.

The cudaProfilerStart and ./pipeline 4 lines are for the next section.

Seeing the overlap with Nsight Systems

Timings only say that the pipeline got faster. A trace shows what ran at the same time. nsys can record every copy and kernel; --capture-range=cudaProfilerApi records only what happens between cudaProfilerStart and cudaProfilerStop, so the warm-up run is left out. This script turns the trace’s CSV into a timeline and adds up how long a copy and a kernel ran together.

# trace.py - read nsys's CUDA GPU trace (CSV) and show what ran when, on which stream
import csv
import sys


def merged(spans):  # join overlapping (start, end) spans into one list of busy periods
    out = []
    for s, e in sorted(spans):
        if out and s <= out[-1][1]:
            out[-1][1] = max(out[-1][1], e)
        else:
            out.append([s, e])
    return out


def shared_time(a, b):  # how long both lists of busy periods are busy at once
    return sum(max(0, min(e1, e2) - max(s1, s2)) for s1, e1 in a for s2, e2 in b)


rows = [r for r in csv.DictReader(sys.stdin) if r.get("Duration (ns)")]
t0 = min(int(r["Start (ns)"]) for r in rows)
copies, kernels = [], []
for r in rows:
    start = (int(r["Start (ns)"]) - t0) / 1e6
    end = start + int(r["Duration (ns)"]) / 1e6
    name = r["Name"]
    kind = "copy in" if "Host-to-Device" in name else "copy out" if "Device-to-Host" in name else "kernel"
    (kernels if kind == "kernel" else copies).append((start, end))
    if "-v" in sys.argv:
        print(f"{start:6.2f} -> {end:6.2f} ms  stream {r['Strm']}  {kind}")
span = max(e for _, e in copies + kernels)
print(f"first operation to last: {span:.2f} ms")
print(f"a copy and a kernel running at the same time: {shared_time(merged(copies), merged(kernels)):.2f} ms")

$ nsys profile --capture-range=cudaProfilerApi -o p1 --force-overwrite true ./pipeline 1 > /dev/null 2>&1
$ nsys stats -q --force-export=true -r cuda_gpu_trace -f csv p1.nsys-rep | python3 trace.py
first operation to last: 18.61 ms
a copy and a kernel running at the same time: 0.00 ms
$ nsys profile --capture-range=cudaProfilerApi -o p4 --force-overwrite true ./pipeline 4 > /dev/null 2>&1
$ nsys stats -q --force-export=true -r cuda_gpu_trace -f csv p4.nsys-rep | python3 trace.py -v
  0.00 ->   0.74 ms  stream 13  copy in
  0.74 ->   1.46 ms  stream 14  copy in
  0.75 ->   1.99 ms  stream 13  kernel
  1.46 ->   2.23 ms  stream 15  copy in
  1.96 ->   3.22 ms  stream 14  kernel
  1.99 ->   2.85 ms  stream 13  copy out
  2.23 ->   3.04 ms  stream 16  copy in
  3.04 ->   3.86 ms  stream 13  copy in
  3.20 ->   4.45 ms  stream 15  kernel
  3.23 ->   4.03 ms  stream 14  copy out
  4.04 ->   4.79 ms  stream 14  copy in
  4.43 ->   5.69 ms  stream 16  kernel
  4.46 ->   5.19 ms  stream 15  copy out
  5.19 ->   5.92 ms  stream 15  copy in
  5.66 ->   6.91 ms  stream 13  kernel
  5.69 ->   6.39 ms  stream 16  copy out
  6.41 ->   7.13 ms  stream 16  copy in
  6.89 ->   8.15 ms  stream 14  kernel
  6.92 ->   7.61 ms  stream 13  copy out
  8.12 ->   9.38 ms  stream 15  kernel
  8.15 ->   8.79 ms  stream 14  copy out
  9.35 ->  10.61 ms  stream 16  kernel
  9.38 ->  10.02 ms  stream 15  copy out
 10.61 ->  11.25 ms  stream 16  copy out
first operation to last: 11.25 ms
a copy and a kernel running at the same time: 8.11 ms

On one stream nothing overlapped: 0.00 ms. On four, a copy ran beside a kernel for 8.11 of the 11.25 ms (four traces: 17.7 to 19.9 ms and no overlap on one stream; 10.1 to 12.2 ms with 7.5 to 8.7 ms of overlap on four). The numbers 13 to 16 are the IDs nsys gave the four streams.

Read the first lines of the four-stream timeline. Stream 13’s kernel starts at 0.75 ms, right after its copy in, while stream 14’s copy in runs from 0.74 to 1.46 ms. Copies in run one at a time, never two at once, because one engine does them all; copies out run on another engine at the same time as the next kernel. The kernels themselves run almost back to back, overlapping only at their tails: each kernel’s 8,192 blocks of 256 threads fill the GPU, so there’s no room for a second kernel. The kernel stage is the floor, which is why four streams couldn’t do much better than two.

The traced run is slower than the timed one: profiling adds work, and the T4’s clock speed varies from run to run. Compare traces with traces, not with the timings above.

The legacy default stream serialises everything

A launch without a stream argument goes to stream 0, the default stream. Here three small kernels, one block each, are launched on stream a, stream 0 and stream b.

// default_stream.cu - one launch on stream 0 makes two other streams wait for each other
#include <chrono>
#include <cstdio>
#include <cuda_runtime.h>

const int ITERS = 1 << 23;

__global__ void spin(float* x) {  // one small block that just burns time
  float v = x[threadIdx.x];
  for (int k = 0; k < ITERS; k++) v = v * 0.5f + 1.0f;
  x[threadIdx.x] = v;
}

double timed(cudaStream_t a, cudaStream_t b, float* x, bool use_stream_0) {
  cudaDeviceSynchronize();
  auto t0 = std::chrono::steady_clock::now();
  spin<<<1, 32, 0, a>>>(x);
  if (use_stream_0) spin<<<1, 32>>>(x + 32);  // no stream argument: the default stream
  spin<<<1, 32, 0, b>>>(x + 64);
  cudaDeviceSynchronize();
  return std::chrono::duration<double, std::milli>(std::chrono::steady_clock::now() - t0).count();
}

int main() {
  float* x; cudaMalloc(&x, 96 * 4); cudaMemset(x, 0, 96 * 4);
  cudaStream_t a, b, na, nb;
  cudaStreamCreate(&a); cudaStreamCreate(&b);
  cudaStreamCreateWithFlags(&na, cudaStreamNonBlocking);
  cudaStreamCreateWithFlags(&nb, cudaStreamNonBlocking);
  for (int i = 0; i < 3; i++) timed(a, b, x, true);  // warm-up: lets the GPU clock ramp up
  printf("streams a and b only:              %6.1f ms\n", timed(a, b, x, false));
  printf("a, then stream 0, then b:          %6.1f ms\n", timed(a, b, x, true));
  printf("same, a and b made NonBlocking:    %6.1f ms\n", timed(na, nb, x, true));
}

$ nvcc -o default_stream default_stream.cu && ./default_stream
streams a and b only:                21.5 ms
a, then stream 0, then b:            64.4 ms
same, a and b made NonBlocking:      21.5 ms
$ nvcc --default-stream per-thread -o default_stream_pt default_stream.cu && ./default_stream_pt
streams a and b only:                21.5 ms
a, then stream 0, then b:            21.5 ms
same, a and b made NonBlocking:      21.5 ms

Each kernel runs for about 21.5 ms, and kernels on a and b alone run side by side. Adding one launch on stream 0 tripled the time to 64.4 ms (64.3 to 64.4 in three runs): the three kernels ran one after another.

By default stream 0 is the legacy default stream, and it synchronises with every stream made by cudaStreamCreate. Work issued to it waits for everything already queued in those streams, and work issued to them afterwards waits for it. So b’s kernel waited for stream 0’s kernel, which waited for a’s. The kernels here use one block each on purpose; overlap needs free room on the GPU, which the small kernels leave.

There are two fixes, and both brought the time back to 21.5 ms:

• cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking): the stream doesn’t synchronise with the legacy default stream at all.
• nvcc --default-stream per-thread: stream 0 becomes a per-thread default stream that behaves like any other stream, for every launch and copy in the files compiled with the flag. A library built without it still uses the legacy default stream unless you pass it a stream.

Stream 0 is easy to use without noticing: cudaMemcpy, cudaMemset and a launch with no stream argument all go there, and so do library calls you haven’t given a stream.

Making one stream wait for another with an event

Sometimes one stream really must wait: a kernel can’t read data that another stream is still copying in. Stream a copies 256 MB of 1.0s into d_in, and stream b adds 1 to every element, first with no dependency between them, then with an event.

// wait_event.cu - stream b reads data that stream a is still copying in, then the fix
#include <cstdio>
#include <cstring>
#include <cuda_runtime.h>

__global__ void add_one(const float* in, float* out, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) out[i] = in[i] + 1.0f;
}

int main(int argc, char** argv) {
  const int n = 1 << 26;  // 256 MB, so the copy in takes about 20 ms
  float *h_in, *h_out, *d_in, *d_out;
  cudaMallocHost(&h_in, n * 4); cudaMallocHost(&h_out, n * 4);
  cudaMalloc(&d_in, n * 4); cudaMalloc(&d_out, n * 4);
  for (int i = 0; i < n; i++) h_in[i] = 1.0f;  // so every output should be 2.0
  cudaStream_t a, b;
  cudaStreamCreateWithFlags(&a, cudaStreamNonBlocking);
  cudaStreamCreateWithFlags(&b, cudaStreamNonBlocking);
  cudaEvent_t copied;
  cudaEventCreateWithFlags(&copied, cudaEventDisableTiming);
  if (argc < 2 || strcmp(argv[1], "nowarm") != 0)  // ./wait_event nowarm skips the warm-up
    add_one<<<1, 1, 0, b>>>(d_in, d_out, 1);       // load the kernel now (see "the race that hid")

  for (int fixed = 0; fixed < 2; fixed++) {
    cudaMemset(d_in, 0, n * 4);
    cudaDeviceSynchronize();
    cudaMemcpyAsync(d_in, h_in, n * 4, cudaMemcpyHostToDevice, a);
    if (fixed) {
      cudaEventRecord(copied, a);         // a marker after the copy in stream a
      cudaStreamWaitEvent(b, copied, 0);  // b's later work waits for it; the CPU doesn't
    }
    add_one<<<n / 256, 256, 0, b>>>(d_in, d_out, n);
    cudaMemcpyAsync(h_out, d_out, n * 4, cudaMemcpyDeviceToHost, b);
    cudaStreamSynchronize(b);
    int wrong = 0;
    for (int i = 0; i < n; i++) wrong += h_out[i] != 2.0f;
    printf("%-26s %8d of %d elements wrong\n", fixed ? "with cudaStreamWaitEvent:" : "no dependency:", wrong, n);
  }
}

$ nvcc -o wait_event wait_event.cu && ./wait_event
no dependency:             67099520 of 67108864 elements wrong
with cudaStreamWaitEvent:         0 of 67108864 elements wrong

Without a dependency, 99.99% of the outputs were wrong (67,098,240 to 67,100,032 in three runs): b’s kernel read d_in while a’s 21.7 ms copy had barely started, so it saw the zeros left by cudaMemset and wrote 1.0. Two streams make no promise about order between them.

cudaEventRecord(copied, a) puts a marker in stream a after the copy, and cudaStreamWaitEvent(b, copied, 0) makes everything queued on b after it wait until the marker is reached. Both calls return to the CPU at once, and stream a and any other stream keep running. With them, 0 elements were wrong. cudaDeviceSynchronize() between the copy and the kernel would also have been correct, but it stops the CPU and waits for every stream on the GPU.

The race that hid

The warm-up launch before the loop is there for a reason. Run the same program with nowarm, which skips it, then again with every kernel loaded at startup.

$ ./wait_event nowarm
no dependency:                    0 of 67108864 elements wrong
with cudaStreamWaitEvent:         0 of 67108864 elements wrong
$ CUDA_MODULE_LOADING=EAGER ./wait_event nowarm
no dependency:             66999168 of 67108864 elements wrong
with cudaStreamWaitEvent:         0 of 67108864 elements wrong

Without the warm-up, the program with no dependency came out right in all three runs. Recent CUDA versions load each kernel lazily, the first time it’s launched. NVIDIA’s lazy-loading notes warn that loading a kernel can synchronise, and here it waited for the copy already running, which happened to put the kernel after it. With CUDA_MODULE_LOADING=EAGER, which loads every kernel at startup, the race came back: 66,964,224 to 67,051,200 elements wrong in three runs.

A missing dependency can pass every test that runs a kernel for the first time. Run stream code with CUDA_MODULE_LOADING=EAGER, or warm up each kernel, before you trust that it’s correct.

Choosing how to wait

Wait for as little as you need, as rarely as you can. Practice 4 measures what waiting too often costs.

Common gotchas

• cudaMemcpyAsync from pageable memory isn’t asynchronous. The 256 MB copy above held the CPU for about 58 ms. Pin the buffers you copy from and to.
• Stream 0 is a synchronisation point. One cudaMemcpy, cudaMemset or launch with no stream serialises every stream made with cudaStreamCreate. Use cudaStreamNonBlocking or --default-stream per-thread.
• Streams don’t make a big kernel share the GPU. Kernels overlap only when the first leaves free room. The pipeline’s kernels each filled the T4 and ran one after another.
• The first launch can hide a race. Lazy loading made the missing dependency above invisible until CUDA_MODULE_LOADING=EAGER exposed it.
• Don’t touch a buffer an async copy is still using. Writing new data into a pinned host buffer before its copy has finished changes what gets copied. Wait on the stream or an event first.
• The first run is slower. With a single warm-up round, the first timing in default_stream.cu was sometimes 33 ms instead of 21.5 ms while the GPU clock ramped up; three rounds fixed it. Time several runs and warm up first.
• CUDA graphs can capture a whole stream pipeline and replay it with one call; see Practice 2 of CUDA performance and profiling. The Python version of this lesson’s pipeline is CUDA Python: streams and overlap.

Practice

1. Keep 4 streams but vary the number of chunks from 1 to 4096. Where does chunking stop helping?

// p1_chunks.cu - the 4-stream pipeline with 1 to 4096 chunks: where does chunking stop paying?
#include <chrono>
#include <cstdio>
#include <cstring>
#include <cuda_runtime.h>

__global__ void work(float* x, int n, int iters) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) {
    float v = x[i];
    for (int k = 0; k < iters; k++) v = v * 0.5f + 1.0f;
    x[i] = v + x[i];
  }
}

const int N = 1 << 24;

double run(int chunks, float* in, float* out, float* d, cudaStream_t* s) {
  int c_len = N / chunks;
  memset(out, 0, N * 4);
  auto t0 = std::chrono::steady_clock::now();
  for (int c = 0; c < chunks; c++) {
    cudaStream_t st = s[c % 4];
    size_t off = (size_t)c * c_len;
    cudaMemcpyAsync(d + off, in + off, c_len * 4, cudaMemcpyHostToDevice, st);
    work<<<(c_len + 255) / 256, 256, 0, st>>>(d + off, c_len, 1000);
    cudaMemcpyAsync(out + off, d + off, c_len * 4, cudaMemcpyDeviceToHost, st);
  }
  cudaDeviceSynchronize();
  return std::chrono::duration<double, std::milli>(std::chrono::steady_clock::now() - t0).count();
}

int main() {
  float *in, *out, *d;
  cudaMallocHost(&in, N * 4); cudaMallocHost(&out, N * 4); cudaMalloc(&d, N * 4);
  for (int i = 0; i < N; i++) in[i] = i % 100;
  cudaStream_t s[4];
  for (int i = 0; i < 4; i++) cudaStreamCreate(&s[i]);
  run(8, in, out, d, s);  // warm-up
  for (int chunks : {1, 2, 4, 8, 16, 32, 64, 256, 1024, 4096}) {
    double best = 1e9;
    for (int rep = 0; rep < 5; rep++) { double t = run(chunks, in, out, d, s); best = t < best ? t : best; }
    int wrong = 0;
    for (int i = 0; i < N; i++) wrong += out[i] != 2.0f + i % 100;
    printf("%5d chunks of %8d floats: %6.2f ms (best of 5), wrong: %d\n", chunks, N / chunks, best, wrong);
  }
}

$ nvcc -o p1_chunks p1_chunks.cu && ./p1_chunks
    1 chunks of 16777216 floats:  16.96 ms (best of 5), wrong: 0
    2 chunks of  8388608 floats:  11.59 ms (best of 5), wrong: 0
    4 chunks of  4194304 floats:   8.97 ms (best of 5), wrong: 0
    8 chunks of  2097152 floats:   7.91 ms (best of 5), wrong: 0
   16 chunks of  1048576 floats:   7.44 ms (best of 5), wrong: 0
   32 chunks of   524288 floats:   7.39 ms (best of 5), wrong: 0
   64 chunks of   262144 floats:   7.51 ms (best of 5), wrong: 0
  256 chunks of    65536 floats:   8.36 ms (best of 5), wrong: 0
 1024 chunks of    16384 floats:  13.44 ms (best of 5), wrong: 0
 4096 chunks of     4096 floats:  42.28 ms (best of 5), wrong: 0

2. Allocate a third buffer with cudaHostAlloc(..., cudaHostAllocWriteCombined) and compare its copy speed to the GPU, and how fast the CPU can read it, with malloc and cudaMallocHost.

// p2_wc.cu - host-to-device copies from pageable, pinned and write-combined pinned memory
#include <chrono>
#include <cstdio>
#include <cstdlib>
#include <cuda_runtime.h>

float h2d_ms(float* h, float* d, size_t bytes) {
  cudaEvent_t a, b; cudaEventCreate(&a); cudaEventCreate(&b);
  float best = 1e9f;
  for (int rep = 0; rep < 5; rep++) {
    cudaEventRecord(a);
    cudaMemcpyAsync(d, h, bytes, cudaMemcpyHostToDevice);
    cudaEventRecord(b);
    cudaEventSynchronize(b);
    float ms; cudaEventElapsedTime(&ms, a, b);
    best = ms < best ? ms : best;
  }
  return best;
}

double cpu_read_ms(const float* h, size_t n) {  // the CPU reads every 16th float (one per 64-byte line)
  auto t0 = std::chrono::steady_clock::now();
  volatile double sum = 0;
  for (size_t i = 0; i < n; i += 16) sum = sum + h[i];
  return std::chrono::duration<double, std::milli>(std::chrono::steady_clock::now() - t0).count();
}

int main() {
  const size_t n = 1 << 26, bytes = n * 4;  // 256 MB
  float *pageable = (float*)malloc(bytes), *pinned, *wc, *d;
  cudaMallocHost(&pinned, bytes);
  cudaHostAlloc(&wc, bytes, cudaHostAllocWriteCombined);
  cudaMalloc(&d, bytes);
  for (size_t i = 0; i < n; i++) pageable[i] = pinned[i] = wc[i] = 1.0f;
  const char* names[] = {"malloc", "cudaMallocHost", "WriteCombined"};
  float* bufs[] = {pageable, pinned, wc};
  for (int k = 0; k < 3; k++) {
    float ms = h2d_ms(bufs[k], d, bytes);
    printf("%-15s copy to GPU %6.2f ms (%5.2f GB/s), CPU read %8.2f ms\n",
           names[k], ms, bytes / ms / 1e6, cpu_read_ms(bufs[k], n));
  }
}

$ nvcc -o p2_wc p2_wc.cu && ./p2_wc
malloc          copy to GPU  57.53 ms ( 4.67 GB/s), CPU read    22.94 ms
cudaMallocHost  copy to GPU  21.71 ms (12.36 GB/s), CPU read    23.58 ms
WriteCombined   copy to GPU  21.69 ms (12.38 GB/s), CPU read    22.51 ms

3. Run a short urgent kernel behind a large one, first on a stream of the same priority, then on a stream created with cudaStreamCreateWithPriority at the highest priority.

// p3_priority.cu - a short urgent kernel behind a big busy one: does a high-priority stream help?
#include <cstdio>
#include <cuda_runtime.h>

__global__ void work(float* x, int n, int iters) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) { float v = x[i]; for (int k = 0; k < iters; k++) v = v * 0.5f + 1.0f; x[i] = v; }
}

float urgent_latency(cudaStream_t big, cudaStream_t urgent, float* x, float* y) {
  cudaEvent_t a, b; cudaEventCreate(&a); cudaEventCreate(&b);
  work<<<1 << 16, 256, 0, big>>>(x, 1 << 24, 4000);  // fills every SM for a while
  cudaEventRecord(a, urgent);
  work<<<4, 256, 0, urgent>>>(y, 1024, 100);           // tiny, launched right after
  cudaEventRecord(b, urgent);
  cudaDeviceSynchronize();
  float ms; cudaEventElapsedTime(&ms, a, b);
  return ms;
}

int main() {
  int least, greatest;
  cudaDeviceGetStreamPriorityRange(&least, &greatest);
  printf("priority range: %d (lowest) to %d (highest)\n", least, greatest);
  float *x, *y; cudaMalloc(&x, (1 << 24) * 4); cudaMalloc(&y, 1024 * 4);
  cudaStream_t big, normal, high;
  cudaStreamCreateWithPriority(&big, cudaStreamNonBlocking, least);
  cudaStreamCreateWithPriority(&normal, cudaStreamNonBlocking, least);
  cudaStreamCreateWithPriority(&high, cudaStreamNonBlocking, greatest);
  urgent_latency(big, normal, x, y);  // warm-up
  for (int rep = 0; rep < 3; rep++) {
    float same = urgent_latency(big, normal, x, y);
    float higher = urgent_latency(big, high, x, y);
    printf("urgent kernel finished after: same priority %6.2f ms, high priority %6.2f ms\n", same, higher);
  }
}

$ nvcc -o p3_priority p3_priority.cu && ./p3_priority
priority range: 0 (lowest) to -5 (highest)
urgent kernel finished after: same priority  16.43 ms, high priority   0.03 ms
urgent kernel finished after: same priority  16.42 ms, high priority   0.03 ms
urgent kernel finished after: same priority  16.54 ms, high priority   0.03 ms

4. How often you wait matters. Put a cudaDeviceSynchronize() after each chunk of the 4-stream pipeline, “just to be safe”, and compare it with waiting once at the end, on one event per stream.

// p4_sync.cu - the 4-stream pipeline with a cudaDeviceSynchronize after every chunk, then one event per stream
#include <chrono>
#include <cstdio>
#include <cstring>
#include <cuda_runtime.h>

__global__ void work(float* x, int n, int iters) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) {
    float v = x[i];
    for (int k = 0; k < iters; k++) v = v * 0.5f + 1.0f;
    x[i] = v + x[i];
  }
}

const int N = 1 << 24, CHUNKS = 8, C = N / CHUNKS;

double run(bool sync_every_chunk, float* in, float* out, float* d, cudaStream_t* s, cudaEvent_t* done) {
  memset(out, 0, N * 4);
  auto t0 = std::chrono::steady_clock::now();
  for (int c = 0; c < CHUNKS; c++) {
    cudaStream_t st = s[c % 4];
    size_t off = (size_t)c * C;
    cudaMemcpyAsync(d + off, in + off, C * 4, cudaMemcpyHostToDevice, st);
    work<<<C / 256, 256, 0, st>>>(d + off, C, 1000);
    cudaMemcpyAsync(out + off, d + off, C * 4, cudaMemcpyDeviceToHost, st);
    if (sync_every_chunk) cudaDeviceSynchronize();  // "just to be safe"
  }
  if (!sync_every_chunk) {
    for (int i = 0; i < 4; i++) cudaEventRecord(done[i], s[i]);
    for (int i = 0; i < 4; i++) cudaEventSynchronize(done[i]);  // wait for each stream's last chunk
  }
  return std::chrono::duration<double, std::milli>(std::chrono::steady_clock::now() - t0).count();
}

int main() {
  float *in, *out, *d;
  cudaMallocHost(&in, N * 4); cudaMallocHost(&out, N * 4); cudaMalloc(&d, N * 4);
  for (int i = 0; i < N; i++) in[i] = i % 100;
  cudaStream_t s[4]; cudaEvent_t done[4];
  for (int i = 0; i < 4; i++) { cudaStreamCreate(&s[i]); cudaEventCreate(&done[i]); }
  run(false, in, out, d, s, done);  // warm-up
  for (bool sync : {true, false}) {
    double total = 0;
    for (int rep = 0; rep < 5; rep++) total += run(sync, in, out, d, s, done);
    int wrong = 0;
    for (int i = 0; i < N; i++) wrong += out[i] != 2.0f + i % 100;
    printf("%-32s %6.2f ms (mean of 5), wrong: %d\n",
           sync ? "cudaDeviceSynchronize per chunk:" : "one event per stream at the end:", total / 5, wrong);
  }
}

$ nvcc -o p4_sync p4_sync.cu && ./p4_sync
cudaDeviceSynchronize per chunk:  15.03 ms (mean of 5), wrong: 0
one event per stream at the end:   7.73 ms (mean of 5), wrong: 0