100 DAYS OF GPU Course preview
Day 53Module 6
in-technical-review

Pinned memory and async copies

Day 51 taught you that copies can overlap kernels, so you did the obvious thing: changed cudaMemcpy to cudaMemcpyAsync, passed a stream, and captured a timeline. Nothing overlaps.

The documentation warns that this can happen: for pageable memory the driver "may synchronize with the stream" before the copy runs. Whether this driver makes the host sit inside the call for most of the copy is exactly what prediction 2 below asks the run to decide.

Either way, the call and stream work as designed. The source memory prevents overlap.

By the end of this page you will know why an ordinary malloc or std::vector buffer cannot be copied asynchronously, what cudaMallocHost changes, and what it costs to skip the copy entirely.

Where a copy actually comes from

The GPU does not read your buffer with a loop. A copy engine on the card performs DMA: it is given a physical address and it pulls the bytes across PCIe with no CPU involved. That only works if the bytes hold still.

An ordinary allocation is pageable, which means the OS may move it, or write it out and reuse the page, at any moment. A physical address may then change, so the driver cannot DMA from it directly.

What the driver does instead is stage. It keeps a modest pinned bounce buffer of its own, copies your pageable bytes into it with the CPU, lets the copy engine DMA out of that, and repeats until it has copied the full 256 MiB. Every byte moves twice and the CPU touches all of them.

NVIDIA states the consequence plainly for the blocking copy: "For transfers from pageable host memory to device memory, a stream sync is performed before the copy is initiated. The function will return once the pageable buffer has been copied to the staging memory for DMA transfer to device memory, but the DMA to final destination may not have completed" (https://docs.nvidia.com/cuda/cuda-runtime-api/api-sync-behavior.html , checked 2026-09-01).

For cudaMemcpyAsync the same page says: "If pageable memory must first be staged to pinned memory, the driver may synchronize with the stream and stage the copy into pinned memory."

cudaMallocHost allocates memory that is page-locked from the start. The OS promises not to move it, the physical address stays true, and the copy engine can DMA straight from your buffer: one crossing, no staging, no CPU in the loop.

The CUDA C++ Best Practices Guide gives this summary: "Page-locked or pinned memory transfers attain the highest bandwidth between the host and the device" (https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html section 10.1.1, checked 2026-09-01).

Diagram: one 64 MiB copy, three ways. Three horizontal bands. Each shows a host buffer on the left, the driver and PCIe in the middle, device memory on the right, arrows for bytes. Band 1, pageable: two arrows in series, CPU from buffer to a small staging box, then DMA from staging to device. Caption "every byte crosses twice, and the CPU carries it halfway." Band 2, pinned: one DMA arrow from buffer to device, with the API call marked as already returned at the arrow's start. Caption "one crossing, no CPU, the call returns before the bytes do." Band 3, zero-copy: no copy arrow; fourteen thin arrows from the device side reaching back into the host buffer, one per kernel launch. Caption "no copy, but 14 launches mean 14 crossings of the same bytes." Alt text: "A pageable copy moves every byte twice through a staging buffer, a pinned copy moves it once by DMA, and a zero-copy kernel re-crosses PCIe on all fourteen launches."

"Async" is a property of the memory, not the call

The function name alone does not make a copy asynchronous. cudaMemcpyAsync is a request, and the runtime grants it only when the source can be DMA'd. The Best Practices Guide means this when it says the asynchronous transfer version "requires pinned host memory" (section 10.1.2, checked 2026-09-01).

Pageable memory still works and returns cudaSuccess, but the driver uses the staged copy above while the host may wait. The API reports no error or warning. Use the timeline to check whether the copy overlapped other work.

Pinning cuts the other way too. Once the copy really is asynchronous, the call returning no longer means the data has been read. Overwrite a pinned source buffer right after cudaMemcpyAsync returns and you are racing the copy engine; whichever bytes it reaches last are the ones the device gets.

The synchronisation you were freed from is now your job, which is what day 52's events are for.

One binary, three questions

Full program in code/day53-pinned/pinned.cu. Three things keep its answers clean.

Identical bytes on every path. The pageable and pinned buffers are filled with the same values, and both paths are gated on a full copy-back comparison before anything is timed, so a fast row cannot be a copy that did not happen. Part 1 then times blocking H2D cudaMemcpy from each source at 1, 4, 16, 64 and 256 MiB, with CUDA events around ten copies after three warm-ups.

The host clock only ever reads settled state. Part 2 issues the same 64 MiB cudaMemcpyAsync on a non-default stream and asks two questions per source: how long did the call hold the thread, and how long until the copy completed. That first question is about host blocking, which no device-side event can see, so this is the course's second sanctioned use of day 9's host clock. The clock is read after the call returns and after cudaStreamSynchronize, never beside in-flight work.

    for (int s = 0; s < 2; ++s) {
        CUDA_CHECK(cudaMemcpyAsync(d_buf, srcs[s], kProbeBytes,
                                   cudaMemcpyHostToDevice, stream));
        CUDA_CHECK(cudaStreamSynchronize(stream));  // warm-up, not timed

        char label[64];
        std::snprintf(label, sizeof(label), "async-%s", names[s]);
        nvtxRangePushA(label);
        double callSum = 0.0;
        double totalSum = 0.0;
        for (int r = 0; r < kTimedRuns; ++r) {
            const double t0 = hostMs();
            CUDA_CHECK(cudaMemcpyAsync(d_buf, srcs[s], kProbeBytes,
                                       cudaMemcpyHostToDevice, stream));
            // The call has returned; how long did it hold the thread?
            const double t1 = hostMs();
            CUDA_CHECK(cudaStreamSynchronize(stream));
            // The copy is done; only now may the clock judge the copy.
            const double t2 = hostMs();
            callSum += t1 - t0;
            totalSum += t2 - t0;
        }
        nvtxRangePop();
        const double callMs = callSum / kTimedRuns;
        const double totalMs = totalSum / kTimedRuns;
        std::printf(
            "%9s  call %8.3f ms  completion %8.3f ms  "
            "call is %5.1f%%\n",
            names[s], callMs, totalMs, 100.0 * callMs / totalMs);
    }

Zero-copy pays per access, so the program makes it pay repeatedly. Part 3 maps a pinned buffer into the device's address space and hands the kernel a pointer straight into host memory:

    float* h_mapped = nullptr;
    float* d_mapped = nullptr;
    CUDA_CHECK(cudaHostAlloc(&h_mapped, kernelBytes, cudaHostAllocMapped));
    CUDA_CHECK(cudaHostGetDevicePointer(&d_mapped, h_mapped, 0));

The same scaleCopy kernel then runs fourteen times over device-resident memory and fourteen times through d_mapped (one correctness gate, three warm-ups, ten timed, per variant). Mapped memory is not cached on the GPU, so every launch re-reads 64 MiB across PCIe, while the resident variant paid PCIe once, in its staging copy. The kernel's reads are fully coalesced, which is the best case for zero-copy.

A strided pattern through a mapped pointer would transfer wasted sectors, as day 11 explains. Each phase has an NVTX range, so the capture has six named groups.

Results

Re-verified 2026-09-02 on the same Tesla T4 with driver 580.173.02 and CUDA 13.0 (V13.0.88). Every ordering and correctness result reproduced. The mid-size pageable copies improved materially, but the 256 MiB ratio barely moved (2.81x to 2.72x), the async call shares remained 99.4 and 0.1 percent, and zero-copy remained about 9.9x slower than device-resident.

Both transcripts and the fresh CUDA 13 report are listed in front matter. The published CUDA 12.6 capture ships with its CSV exports at content/profiles/pinned-memory/: day53-pinned.nsys-rep with day53-pinned_cuda_api_sum.csv, day53-pinned_cuda_api_trace.csv, day53-pinned_cuda_gpu_mem_time_sum.csv and day53-pinned_nvtx_sum.csv beside it.

Size pageable (ms, GB/s) pinned (ms, GB/s) ratio
1 MiB 0.300, 3.5 0.098, 10.7 3.08x
4 MiB 0.926, 4.5 0.359, 11.7 2.58x
16 MiB 3.461, 4.8 1.385, 12.1 2.50x
64 MiB 14.692, 4.6 5.462, 12.3 2.69x
256 MiB 59.176, 4.5 21.747, 12.3 2.72x
Source call (ms) completion (ms) call share
pageable 14.772 14.860 99.4 percent
pinned 0.006 5.461 0.1 percent
Variant ms per launch effective GB/s
device-resident 0.548 245.1
zero-copy 5.431 24.7

The expected shape of the numbers

PCIe links differ across systems, so this lesson predicts orderings and rough ratios instead of one bandwidth value. The run sets the absolute numbers.

  1. Pinned beats pageable at every size, and the gap grows with size. At 256 MiB I expect pageable to cost at least 1.3 times pinned. If the two are within 5 percent everywhere, the driver is staging far better than its own documentation implies, and the lesson's premise dies.
  2. The pageable cudaMemcpyAsync call blocks for most of its copy: call time at least 80 percent of completion time. The pinned call blocks for almost none of it: under 10 percent. This is the pair of rows that tests the page's main claim.
  3. The zero-copy kernel loses to the device-resident one by roughly the ratio of part 1's pinned H2D bandwidth to the device's own bandwidth, within a factor of two. The resident kernel should sit near the card's measured copy bandwidth; the mapped kernel is a PCIe consumer even though a kernel issues the reads.
  4. In the Nsight Systems capture's cuda_api_trace export, the ten timed cudaMemcpyAsync calls from the pageable source cost at least five times more total call time than the ten from the pinned source, even though cuda_gpu_mem_time_sum shows both moving the same bytes. The trace makes the two groups findable with no GUI open: part 2's copies are the program's only cudaMemcpyAsync calls, eleven per source (one warm-up, then the ten timed), with the pageable eleven first.

Prediction 2 tests the driver behavior directly. If the pageable call returns quickly and the copy still completes late, the driver is staging asynchronously. The result would support "pageable async is slow" rather than "pageable async blocks the host".

Those cases need different fixes, and the timeline distinguishes them.

What the run said

  1. The floor held; the growth claim died again. Pageable cost 2.72x pinned at 256 MiB, past the 1.3x bar. But the gap does not grow with size: it wanders (3.08, 2.58, 2.50, 2.69, 2.72). Staging cost is not a simple per-byte tax on this driver.
  2. Held, emphatically, and it is the page's title. The pageable call held the thread for 99.4 percent of its copy's completion time; the pinned call for 0.1 percent. On this driver, pageable async really is secretly synchronous.
  3. Died by a hair, in an instructive direction. Part 1's pinned rate (12.3 GB/s) against the resident kernel's 245.1 GB/s predicts a 19.9x loss; measured is 9.92x, again just outside the factor-of-two bound. The mapped reads sustained 24.7 GB/s, twice the copy engine's H2D rate: per-lane reads across PCIe pipeline better than the blocking copy path the prediction priced them against.
  4. Held with two orders of magnitude to spare in the published CUDA 12.6 trace. In cuda_api_trace, the ten timed pageable cudaMemcpyAsync calls total 146.556 ms of call time against 0.085 ms for the ten pinned ones, a 1,717x gap against the predicted 5x floor, while cuda_gpu_mem_time_sum shows both groups moving the same bytes.

Use pinned memory for buffers that feed cudaMemcpyAsync. Allocate them with cudaMallocHost, or the copy is unlikely to overlap other work. The tables test that rule.

Run it yourself

Use a CUDA GPU with enough host and device memory for the listed buffers.

nvcc -std=c++17 -O3 -arch=sm_75 -lineinfo -o pinned pinned.cu
./pinned
nsys profile -t cuda,nvtx -o day53-pinned ./pinned

There is no Compiler Explorer embed: the program pins 320 MiB of host memory, which is too much for a shared sandbox, and the artifact that matters is the nsys capture, which no embed can produce. If you have no GPU, read /setup/learn-cuda-without-a-gpu and work from the shipped capture at content/profiles/pinned-memory/.

Exercise

Add a device-to-host sweep to part 1: the same five sizes, the same two destinations, direction reversed. Report the pageable/pinned ratio at 64 MiB in both directions and say whether pinning matters more, less or the same on the way back.

Time: 25 to 35 minutes. Submit: the two-direction ratio table and one sentence naming the direction where pinning buys more.

Check: ratios, not milliseconds, because a laptop's PCIe and a datacenter's give different absolute times. Your H2D pinned/pageable ratio at 64 MiB should be within 20 percent of the reference table's once the reference exists; the D2H ratio is yours to discover. If either pinned row is slower than its pageable row, your "pinned" buffer probably is not: check that the allocation actually came from cudaMallocHost and that the return value was checked.

Hint 1

The staging story in the mental model was told for host-to-device. Walk it backwards: when the device sends bytes to a pageable buffer, who has to touch them before they land where you asked?

Hint 2

The sync-behavior page gives D2H its own rule: "For transfers from device to either pageable or pinned host memory, the function returns only once the copy has completed." So the blocking behavior stops distinguishing the two buffers in this direction. What is left to differ is the staging itself.

Solution

Pinning wins in both directions because the bounce buffer and the double movement of every byte exist in both. D2H from the device lands in the driver's staging area by DMA and the CPU then copies it out to your pageable buffer, so the CPU touches every byte in either direction.

Expect broadly similar ratios rather than a dramatic asymmetry, and expect both blocking calls to return only at completion, per the rule Hint 2 quotes.

Pinning is a property of pages, not of a direction or an API. Any transfer that ends or begins in memory the OS can move needs staging, whichever way the bytes flow.

Pitfalls

You added Async and a stream and nothing overlaps. The source is pageable, so the driver "may synchronize with the stream and stage the copy into pinned memory", which serialises exactly what you tried to overlap. Allocate the staging side of every async copy with cudaMallocHost. This page's part 2 measures the difference.

You pinned everything and the whole machine got slower. Pinned pages are removed from the OS's working set for the life of the allocation. The Best Practices Guide: "Excessive use can reduce overall system performance because pinned memory is a scarce resource, but how much is too much is difficult to know in advance." Pin transfer buffers, not data sets.

You used zero-copy as a general-purpose way to skip copies. Mapped memory is not cached on the GPU: "mapped pinned memory should be read or written only once, and the global loads and stores that read and write the memory should be coalesced." Touch it twice and you paid PCIe twice. Part 3 makes this a printed ratio.

You overwrote the source buffer right after cudaMemcpyAsync returned. With a pinned source the call returns before the DMA reads your bytes, so the write races the copy engine and the device receives a mixture. Synchronise the stream, or record an event and wait on it, before reusing the buffer. Day 52 covers the tools.

You confused mapped pinned memory with unified memory. Both give the device a pointer to something host-visible, but mapped memory never migrates: every access crosses PCIe, forever. Unified memory moves pages to the touching processor and caches them there. Day 19 measured the migration; this page prices the alternative that never migrates.

Go deeper

Next

Day 54 is where this stops being a microbenchmark: pinned staging buffers plus cudaMemcpyAsync plus two streams let a 1 GB array flow through a 64 MB device window, with the copy engine filling one buffer while the SMs process the other. Every piece of that pipeline is something this page just priced.