The Overlap Benchmark: Serial vs Pipelined

By the end of this page you have a benchmark that pits three ways of running the same work against each other, and you know why the middle one — “async but one stream” — is the trap.

The workload

Copy 64 MiB of floats from host to device, scale every element (a deliberately heavy kernel), copy the 64 MiB back. Repeated across 16 chunks of 4 MiB. Chunking is not optional: overlap needs interleaving — the DMA engine must have a next copy queued while the SMs run a kernel. With one giant copy there is nothing to interleave.

The three variants

Variant	Host memory	Structure	What it isolates
`serial`	pageable (`malloc`)	sync `cudaMemcpy` + kernel, default stream	the “before streams” baseline
`naive`	pinned	chunked `cudaMemcpyAsync`, one stream	async API with no structure
`pipelined`	pinned	3 streams + events: H2D copies, kernels, D2H copies	the real thing

The naive variant is the guide’s key lesson: cudaMemcpyAsync is not overlap. Everything lands in one in-order queue — copy, kernel, copy — so the queue serializes exactly like the serial version. It’s the version people write when they “convert to async” without restructuring, and it buys nothing.

The pipelined variant puts each kind of work in its own stream and chains them with events: chunk i’s copy lands in sH2D, an event is recorded, sKernel waits on it and runs the kernel, another event, sD2H waits and copies back. Because the streams are independent, the DMA engine can be copying chunk i+1 in while the SMs scale chunk i — the copy engine and the SMs are busy simultaneously.

The program

overlap.cu (complete; this exact file compiled clean with nvcc 13.2.86):

// overlap.cu — serial vs naive-async vs pipelined copy/compute overlap
// (CUDA Toolkit 13.2)
//
// Compile:  nvcc -arch=sm_80 -O2 overlap.cu -o overlap
// Run:      ./overlap
//
// Three ways to run the same work (copy in 64 MiB, scale it, copy out):
//   1. serial    — synchronous cudaMemcpy + kernel on the default stream
//   2. naive     — one non-blocking stream, chunked cudaMemcpyAsync
//                  (async API, but everything is in ONE queue: no overlap)
//   3. pipelined — three streams + events: H2D copies, kernels, D2H copies
//                  each in their own stream, chained with events
//
// Tune ITERS so kernel time ~= transfer time on your GPU; that is where
// overlap shows the biggest win.

#include <cstdio>
#include <cstdlib>
#include <cstring>

#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)

#define CHUNKS 16
#define FLOATS_PER_CHUNK (1 << 20)  // 1 Mi floats = 4 MiB per chunk
#define TOTAL_FLOATS (CHUNKS * FLOATS_PER_CHUNK)
#define ITERS 5000                  // tune: more = heavier kernel
#define FACTOR 1.0000001f

__global__ void scale_kernel(float *data, int n, float factor, int iters) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) {
    float v = data[i];
    for (int it = 0; it < iters; ++it) v *= factor;
    data[i] = v;
  }
}

static float expected_value(float x, float factor, int iters) {
  for (int it = 0; it < iters; ++it) x *= factor;
  return x;
}

// Bit-exact check: the device runs the same float ops in the same order,
// so results must match the host-computed reference exactly.
static int verify(const float *got, const float *ref, int n, const char *what) {
  for (int i = 0; i < n; ++i) {
    if (got[i] != ref[i]) {
      fprintf(stderr, "%s: WRONG at %d: got %.8g, want %.8g\n", what, i,
              (double)got[i], (double)ref[i]);
      return 1;
    }
  }
  printf("%-22s results OK\n", what);
  return 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  (asyncEngineCount = %d)\n", prop.name, prop.asyncEngineCount);
  if (prop.asyncEngineCount == 0)
    printf("NOTE: no async copy engine — copy/kernel overlap is not possible on this device\n");

  const size_t chunkBytes = (size_t)FLOATS_PER_CHUNK * sizeof(float);
  const size_t totalBytes = (size_t)TOTAL_FLOATS * sizeof(float);

  // Pinned host buffers (used by the async variants).
  float *h_in_p, *h_out_p;
  CUDA_CHECK(cudaMallocHost(&h_in_p, totalBytes));
  CUDA_CHECK(cudaMallocHost(&h_out_p, totalBytes));

  // Pageable host buffers (the "before streams" serial baseline).
  float *h_in_s = (float *)malloc(totalBytes);
  float *h_out_s = (float *)malloc(totalBytes);

  float *d_in;
  CUDA_CHECK(cudaMalloc(&d_in, totalBytes));

  float *h_ref = (float *)malloc(totalBytes);
  for (int i = 0; i < TOTAL_FLOATS; ++i) {
    h_in_p[i] = 1.0f;
    h_ref[i] = expected_value(1.0f, FACTOR, ITERS);
  }
  memcpy(h_in_s, h_in_p, totalBytes);

  cudaEvent_t t0, t1;
  CUDA_CHECK(cudaEventCreate(&t0));
  CUDA_CHECK(cudaEventCreate(&t1));

  const int grid = FLOATS_PER_CHUNK / 256;
  float serial_ms = 0.0f, naive_ms = 0.0f, pipe_ms = 0.0f;

  // ---- 1. serial: synchronous copies, default stream ----
  {
    CUDA_CHECK(cudaEventRecord(t0));
    CUDA_CHECK(cudaMemcpy(d_in, h_in_s, totalBytes, cudaMemcpyHostToDevice));
    scale_kernel<<<grid, 256>>>(d_in, TOTAL_FLOATS, FACTOR, ITERS);  // default stream
    CUDA_CHECK(cudaMemcpy(h_out_s, d_in, totalBytes, cudaMemcpyDeviceToHost));
    CUDA_CHECK(cudaEventRecord(t1));
    CUDA_CHECK(cudaEventSynchronize(t1));
    CUDA_CHECK(cudaEventElapsedTime(&serial_ms, t0, t1));
    if (verify(h_out_s, h_ref, TOTAL_FLOATS, "serial")) return 1;
  }

  // ---- 2. naive: one stream, chunked async copies (no overlap) ----
  cudaStream_t s1;
  CUDA_CHECK(cudaStreamCreateWithFlags(&s1, cudaStreamNonBlocking));
  {
    CUDA_CHECK(cudaEventRecord(t0));
    for (int i = 0; i < CHUNKS; ++i) {
      size_t off = (size_t)i * FLOATS_PER_CHUNK;
      CUDA_CHECK(cudaMemcpyAsync(d_in + off, h_in_p + off, chunkBytes,
                                 cudaMemcpyHostToDevice, s1));
      scale_kernel<<<grid, 256, 0, s1>>>(d_in + off, FLOATS_PER_CHUNK, FACTOR, ITERS);
      CUDA_CHECK(cudaMemcpyAsync(h_out_p + off, d_in + off, chunkBytes,
                                 cudaMemcpyDeviceToHost, s1));
    }
    CUDA_CHECK(cudaStreamSynchronize(s1));
    CUDA_CHECK(cudaEventRecord(t1));
    CUDA_CHECK(cudaEventSynchronize(t1));
    CUDA_CHECK(cudaEventElapsedTime(&naive_ms, t0, t1));
    if (verify(h_out_p, h_ref, TOTAL_FLOATS, "naive async")) return 1;
  }

  // ---- 3. pipelined: three streams chained with events ----
  cudaStream_t sH2D, sKernel, sD2H;
  CUDA_CHECK(cudaStreamCreateWithFlags(&sH2D, cudaStreamNonBlocking));
  CUDA_CHECK(cudaStreamCreateWithFlags(&sKernel, cudaStreamNonBlocking));
  CUDA_CHECK(cudaStreamCreateWithFlags(&sD2H, cudaStreamNonBlocking));
  cudaEvent_t h2dDone[CHUNKS], kernelDone[CHUNKS];
  for (int i = 0; i < CHUNKS; ++i) {
    CUDA_CHECK(cudaEventCreateWithFlags(&h2dDone[i], cudaEventDisableTiming));
    CUDA_CHECK(cudaEventCreateWithFlags(&kernelDone[i], cudaEventDisableTiming));
  }
  {
    CUDA_CHECK(cudaEventRecord(t0));
    for (int i = 0; i < CHUNKS; ++i) {
      size_t off = (size_t)i * FLOATS_PER_CHUNK;
      CUDA_CHECK(cudaMemcpyAsync(d_in + off, h_in_p + off, chunkBytes,
                                 cudaMemcpyHostToDevice, sH2D));
      CUDA_CHECK(cudaEventRecord(h2dDone[i], sH2D));
      CUDA_CHECK(cudaStreamWaitEvent(sKernel, h2dDone[i], 0));
      scale_kernel<<<grid, 256, 0, sKernel>>>(d_in + off, FLOATS_PER_CHUNK,
                                              FACTOR, ITERS);
      CUDA_CHECK(cudaEventRecord(kernelDone[i], sKernel));
      CUDA_CHECK(cudaStreamWaitEvent(sD2H, kernelDone[i], 0));
      CUDA_CHECK(cudaMemcpyAsync(h_out_p + off, d_in + off, chunkBytes,
                                 cudaMemcpyDeviceToHost, sD2H));
    }
    CUDA_CHECK(cudaDeviceSynchronize());
    CUDA_CHECK(cudaEventRecord(t1));
    CUDA_CHECK(cudaEventSynchronize(t1));
    CUDA_CHECK(cudaEventElapsedTime(&pipe_ms, t0, t1));
    if (verify(h_out_p, h_ref, TOTAL_FLOATS, "pipelined")) return 1;
  }

  printf("\n%-12s %10s %12s\n", "variant", "time (ms)", "vs serial");
  printf("%-12s %10.2f %12.2fx\n", "serial", serial_ms, 1.0);
  printf("%-12s %10.2f %12.2fx\n", "naive", naive_ms, serial_ms / naive_ms);
  printf("%-12s %10.2f %12.2fx\n", "pipelined", pipe_ms, serial_ms / pipe_ms);

  for (int i = 0; i < CHUNKS; ++i) {
    CUDA_CHECK(cudaEventDestroy(h2dDone[i]));
    CUDA_CHECK(cudaEventDestroy(kernelDone[i]));
  }
  CUDA_CHECK(cudaEventDestroy(t0));
  CUDA_CHECK(cudaEventDestroy(t1));
  CUDA_CHECK(cudaStreamDestroy(s1));
  CUDA_CHECK(cudaStreamDestroy(sH2D));
  CUDA_CHECK(cudaStreamDestroy(sKernel));
  CUDA_CHECK(cudaStreamDestroy(sD2H));
  CUDA_CHECK(cudaFree(d_in));
  CUDA_CHECK(cudaFreeHost(h_in_p));
  CUDA_CHECK(cudaFreeHost(h_out_p));
  free(h_in_s);
  free(h_out_s);
  free(h_ref);
  return 0;
}

Compile and run (compile verified; the run is yours):

nvcc -arch=sm_80 -O2 overlap.cu -o overlap
./overlap

What to expect (expectations, not measurements)

This guide’s build host has no GPU, so no benchmark numbers were captured here. The following is what the three variants should do on a discrete GPU with asyncEngineCount > 0 — check it against your run:

serial is the sum of everything: full transfer time + full kernel time. It’s the upper bound.
naive lands at roughly serial’s time (sometimes a few percent better, from pinned memory; sometimes a few percent worse, from chunking overhead). It must not be dramatically faster — if it is, your serial baseline was already paying a pageable-staging penalty and the comparison is contaminated.
pipelined approaches max(transfer time, kernel time) plus a small ramp. With ITERS tuned so kernel time ≈ transfer time, expect a real speedup — on the order of 1.5–2× versus serial — because the DMA engine and the SMs are each busy for the whole run instead of alternating.

Output shape (real program, expected values):

GPU: NVIDIA GeForce RTX 3090  (asyncEngineCount = 2)
serial               results OK
naive async          results OK
pipelined            results OK

variant      time (ms)    vs serial
serial           42.31        1.00x
naive            41.87        1.01x
pipelined        24.15        1.75x

If your machine prints asyncEngineCount = 0, pipelined ≈ serial is the correct outcome — the hardware can’t overlap, and no code structure can change that.

Tuning knobs

ITERS — the main knob. Kernel time should be comparable to chunk transfer time. Too small: the run is transfer-bound and overlap saves little (serial ≈ pipelined). Too large: compute-bound, same story. The sweet spot is where the two engines are balanced, and it moves with your GPU’s SM count and PCIe generation. Try 1000, 5000, 20000.
CHUNKS — more chunks = finer interleaving, but each launch/copy has overhead (~5–10 µs). 8–32 chunks is a sane range for this size; at 100+ chunks the overhead eats the gains.
Chunk size — 1–4 MiB per chunk is typical for PCIe 3.0/4.0 machines; much smaller chunks underutilize the DMA engine’s burst efficiency.

When overlap doesn’t help (and how to tell)

Transfer-bound: kernel is trivial next to the copies. Serial is already ≈ transfer time; overlap can save at most the kernel time. Fix the kernel, or accept it.
No copy engine (asyncEngineCount == 0): nothing to overlap with.
Small transfers: below a few hundred KB, launch and sync overhead dwarfs the transfer. Don’t pipeline 4 KB chunks.
Already saturated: if the bus is the bottleneck in both directions and the kernel is tiny, there’s nothing left to hide.

The cheap diagnostic is Nsight Systems: in the timeline, the pipelined variant should show copy and kernel rows interleaved, not sequential. If the rows are sequential, check pinned memory first, then default-stream usage, then the event chain.

Checklist for your own pipeline

Host buffers involved in async copies are cudaMallocHost, not malloc.
Every stream created with cudaStreamNonBlocking; no stray default-stream ops in the hot path.
Every cross-stream dependency is an explicit cudaStreamWaitEvent — never implicit sync.
Sync-only events use cudaEventDisableTiming; timing events don’t.
Results are verified against a reference, not eyeballed.
Timings are warm, measured with events (not clock()), on the GPU you’re actually shipping to.