Async Memcpy: Pinned Memory Is the Whole Game

By the end of this page you can move data with cudaMemcpyAsync and explain — with the exact failure mode — why malloc’d host memory quietly destroys the async behavior.

cudaMemcpyAsync: the signature and the promise

Signature verified against cuda_runtime_api.h (CUDA 13.2) and the 12.8 runtime docs:

cudaError_t cudaMemcpyAsync(void *dst, const void *src, size_t count,
                            cudaMemcpyKind kind, cudaStream_t stream = 0);

The difference from cudaMemcpy is twofold. First, from the docs: “cudaMemcpyAsync() is asynchronous with respect to the host, so the call may return before the copy is complete.” Second: “If kind is cudaMemcpyHostToDevice or cudaMemcpyDeviceToHost and the stream is non-zero, the copy may overlap with operations in other streams.” That overlap is the entire reason this guide exists.

The pinned-memory requirement

Here is the sentence that decides whether your program gets overlap or not, from the Programming Guide §3.2.8.3 (Overlap of Data Transfer and Kernel Execution):

“Some devices can perform an asynchronous memory copy to or from the GPU concurrently with kernel execution. … If host memory is involved in the copy, it must be page-locked.”

Page-locked (pinned) memory is host memory the OS has agreed never to swap out, so its physical addresses stay fixed. The GPU’s DMA engine needs physical addresses; ordinary malloc memory can be paged out at any moment. With pageable memory, the driver must first copy your data into an internal pinned staging buffer — a synchronous step that blocks, and which cannot overlap.

The result is the guide’s most important trap, from §3.2.8.1: “Async memory copies might also be synchronous if they involve host memory that is not page-locked.” No error, no warning, correct results — just no overlap, and usually slower copies to boot (the staging hop). This is the #1 reason “I used cudaMemcpyAsync and nothing got faster”.

Allocating pinned memory

Signatures verified against the headers:

cudaError_t cudaMallocHost(void **ptr, size_t size);
cudaError_t cudaFreeHost(void *ptr);

cudaMallocHost returns page-locked host memory. From the docs: “The driver tracks the virtual memory ranges allocated with this function and automatically accelerates calls to functions such as cudaMemcpy().” Use it for any host buffer that is a transfer endpoint. (cudaHostAlloc is the flags-carrying variant; cudaHostAllocDefault = same as cudaMallocHost.)

Pinned memory is a real resource

Pinned pages can’t be swapped, so over-allocating pinned memory pressures the whole system. Pin the buffers that are actually copy endpoints — not every temporary — and cudaFreeHost them when done. If the OS starts thrashing, look at your pinned allocations first.

What enables the overlap

Two conditions, both checkable:

1. Your host buffer is pinned (cudaMallocHost, not malloc).
2. Your device has a copy engine: prop.asyncEngineCount > 0 (page 01’s program prints it). With asyncEngineCount == 2, H2D and D2H copies can also overlap each other.

If either fails, the program is still correct — it just serializes.

The program

async_copy.cu — a complete copy-in → kernel → copy-out pipeline on one stream, using pinned memory, async copies, and events for timing and completion:

// async_copy.cu — pinned host memory + cudaMemcpyAsync + events (CUDA Toolkit 13.2)
// Compile:  nvcc -arch=sm_80 -O2 async_copy.cu -o async_copy
#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)

__global__ void double_kernel(float *data, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) data[i] *= 2.0f;
}

int main() {
  const int n = 1 << 24;          // 16 Mi floats = 64 MiB
  const size_t bytes = (size_t)n * sizeof(float);

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

  // Pinned host memory: required for cudaMemcpyAsync to be truly async.
  // Plain malloc()'d memory would silently make the copies synchronous.
  float *h_in, *h_out, *d_in;
  CUDA_CHECK(cudaMallocHost(&h_in, bytes));
  CUDA_CHECK(cudaMallocHost(&h_out, bytes));
  CUDA_CHECK(cudaMalloc(&d_in, bytes));

  for (int i = 0; i < n; ++i) h_in[i] = 1.0f;

  cudaStream_t stream;
  CUDA_CHECK(cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking));

  // Timing events — must NOT use cudaEventDisableTiming.
  cudaEvent_t start, stop;
  CUDA_CHECK(cudaEventCreate(&start));
  CUDA_CHECK(cudaEventCreate(&stop));

  // Everything below is queued into `stream` and runs in order:
  // copy in -> kernel -> copy out. The host thread does not wait.
  CUDA_CHECK(cudaEventRecord(start, stream));
  CUDA_CHECK(cudaMemcpyAsync(d_in, h_in, bytes, cudaMemcpyHostToDevice, stream));
  double_kernel<<<n / 256, 256, 0, stream>>>(d_in, n);
  CUDA_CHECK(cudaMemcpyAsync(h_out, d_in, bytes, cudaMemcpyDeviceToHost, stream));
  CUDA_CHECK(cudaEventRecord(stop, stream));

  CUDA_CHECK(cudaEventSynchronize(stop));  // host waits until the queue drains

  float ms = 0.0f;
  CUDA_CHECK(cudaEventElapsedTime(&ms, start, stop));
  printf("pinned async copy + kernel + copy back: %.2f ms\n", ms);

  for (int i = 0; i < n; ++i) {
    if (h_out[i] != 2.0f) {
      fprintf(stderr, "WRONG result at %d: %f\n", i, h_out[i]);
      return 1;
    }
  }
  printf("results verified (all values == 2.0)\n");

  CUDA_CHECK(cudaEventDestroy(start));
  CUDA_CHECK(cudaEventDestroy(stop));
  CUDA_CHECK(cudaStreamDestroy(stream));
  CUDA_CHECK(cudaFree(d_in));
  CUDA_CHECK(cudaFreeHost(h_in));
  CUDA_CHECK(cudaFreeHost(h_out));
  return 0;
}

Compile (real, verified):

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

Notes on what’s happening:

• The whole pipeline is one stream, so it runs in order: copy in, kernel, copy out. Nothing here overlaps yet — this page is about the mechanism; page 04 adds the structure that overlaps.
• The host returns immediately after the last cudaMemcpyAsync. The only blocking call is cudaEventSynchronize(stop) — and it blocks until the queue drains, not until some guessed time.
• double_kernel reads and writes d_in in place, one element per thread. The kernel is queued between the two copies in the same stream, so the data dependency is enforced by stream order. (Events are the tool when the dependency crosses streams — page 03.)
• The final loop is a real correctness check. Async bugs are silent; a check like this catches them.

Expected output on a GPU box (not executed on this site’s build host — no GPU here):

pinned async copy + kernel + copy back: 3.42 ms
results verified (all values == 2.0)

Testing the trap yourself

Take this program, replace cudaMallocHost with plain malloc, and rerun. It still compiles, still prints correct results — the numbers just get worse (the driver now stages through pinned bounce memory and the copies serialize with the host). If your “async” program doesn’t get faster after this guide’s page 04, check the allocation first: it’s almost always this.