Zero-Copy Memory (Pinned Host), in depth: what a kernel load over PCIe costs, the break-even against copying, sparse gathers and host-visible results

What a zero-copy load costs

When a warp loads from an address that maps to host memory, the GPU's memory system cannot satisfy it from HBM. It issues read requests across the interconnect; on a discrete GPU that means PCIe transactions to the host's root complex, which reads DRAM and returns completions. Round-trip latency is on the order of a microsecond, several times HBM latency, and bandwidth is capped by the link: a PCIe Gen4 x16 link is about 32 GB/s per direction before protocol overhead, Gen5 about double. HBM on a current data-centre GPU delivers terabytes per second. The PCIe article explains where those link numbers come from.

Latency matters less than it seems, because GPUs hide latency with parallelism. By Little's law, the bytes in flight needed to fill a link equal bandwidth times latency: 25 GB/s times 1 microsecond is about 25 KB, roughly two hundred 128-byte requests across the whole GPU. A kernel with reasonable occupancy keeps that many outstanding. What hurts is access shape. Every request carries header overhead, so a warp that reads 32 consecutive 4-byte values in one coalesced 128-byte request uses the link efficiently, while a warp that scatters 32 separate 4-byte reads wastes most of it. Coalescing matters for HBM; over a link it decides whether zero-copy is viable at all.

Mapping host memory: the API

Two calls produce mapped memory. cudaHostAlloc with cudaHostAllocMapped allocates new pinned, mapped pages; cudaHostRegister with cudaHostRegisterMapped pins and maps an existing allocation, which is how you expose a buffer some other library created. cudaHostGetDevicePointer returns the address a kernel should use. The runtime documentation notes that the mapped flag only takes effect if the context supports cudaDeviceMapHost, which you can check with cudaGetDeviceFlags; on 64-bit platforms with unified virtual addressing this is normally the case and the device pointer for cudaHostAlloc memory equals the host pointer. For registered memory the two match only on devices reporting cudaDevAttrCanUseHostPointerForRegisteredMem, so always ask for the device pointer rather than assuming.

#include <cuda_runtime.h>
#include <cstdio>
#define CK(x) do { cudaError_t e = (x); if (e != cudaSuccess) { \
    fprintf(stderr, "%s: %s\n", #x, cudaGetErrorString(e)); return 1; } } while (0)

__global__ void scale(const float* __restrict__ in, float* __restrict__ out,
                      float a, size_t n) {
    size_t i = blockIdx.x * (size_t)blockDim.x + threadIdx.x;
    if (i < n) out[i] = a * in[i];          // one coalesced read, one write per element
}

int main() {
    int dev = 0, canMap = 0;
    CK(cudaDeviceGetAttribute(&canMap, cudaDevAttrCanMapHostMemory, dev));
    if (!canMap) { fprintf(stderr, "device cannot map host memory\n"); return 1; }

    size_t n = 1 << 26;                     // 256 MiB of floats
    float *h_in, *h_out, *d_in, *d_out;
    CK(cudaHostAlloc(&h_in,  n * sizeof(float), cudaHostAllocMapped));
    CK(cudaHostAlloc(&h_out, n * sizeof(float), cudaHostAllocMapped));
    for (size_t i = 0; i < n; ++i) h_in[i] = 1.0f;
    CK(cudaHostGetDevicePointer(&d_in,  h_in,  0));   // do not assume d == h
    CK(cudaHostGetDevicePointer(&d_out, h_out, 0));

    cudaStream_t s; CK(cudaStreamCreate(&s));
    scale<<<(n + 255) / 256, 256, 0, s>>>(d_in, d_out, 2.0f, n);
    CK(cudaGetLastError());
    CK(cudaStreamSynchronize(s));           // the host must not read h_out before this
    printf("%f\n", h_out[n - 1]);
    CK(cudaFreeHost(h_in)); CK(cudaFreeHost(h_out));
    return 0;
}

To expose an existing buffer instead, call cudaHostRegister(ptr, bytes, cudaHostRegisterMapped), then cudaHostGetDevicePointer, and cudaHostUnregister before the owner frees it. Adding cudaHostRegisterReadOnly tells the driver the device will only read the range. Registration is expensive, so do it once at start-up, never per batch.

The break-even and the trade-offs

The decision reduces to arithmetic. Let B be the buffer size, r the number of times the kernel reads each byte, L the effective link bandwidth and H the HBM bandwidth. Copying costs about B/L for the transfer plus r times B/H for the reads. Zero-copy costs about r times B/L, because nothing is cached for reuse that you can rely on. With H tens of times larger than L, copying wins as soon as r is noticeably above 1.

At r equal to 1 the two are close in transfer time, and zero-copy's advantages are structural: no device memory is consumed, no separate copy has to be scheduled, and compute starts on the first bytes instead of after the last. A copy pipeline built with multiple streams recovers most of that overlap, so for dense, streaming access the honest answer is often that both approaches reach link speed and copying is easier to reason about.

Pattern 1: sparse gathers from huge host tables

The clearest win is a sparse gather from a table too large for device memory. Graph neural network training is the canonical case: node features live on the host, and each mini-batch needs the rows of a few thousand sampled nodes. DGL, for example, offers a UVA mode that pins the feature table and lets GPU kernels gather directly from it. The same idea applies to recommendation embedding tables and to any lookup index that does not fit in HBM.

A worked example with estimates: 100 million nodes with 256 float32 features is about 102 GB, far beyond one GPU. A batch of 50,000 nodes needs 50,000 rows of 1 KB, about 51 MB. The CPU path gathers those rows into a pinned staging buffer, which means 50,000 random DRAM reads on a few CPU threads and often takes longer than the copy that follows, about 2 milliseconds for 51 MB at 25 GB/s. The zero-copy path assigns one warp per row; each 1 KB row becomes eight coalesced 128-byte requests, the CPU does nothing, and the transfer runs close to link speed, so the batch arrives in roughly the time of the copy alone. Measure on your own system; the gain depends on how slow the CPU gather was.

// One warp per requested row; each lane copies 4-byte words, so a warp's
// 32 lanes read 128 contiguous bytes per iteration (coalesced over the link).
__global__ void gather_rows(const float* __restrict__ host_table,  // mapped
                            const int64_t* __restrict__ ids, int n_ids,
                            int dim, float* __restrict__ out) {    // HBM
    int warp = (blockIdx.x * blockDim.x + threadIdx.x) / 32;
    int lane = threadIdx.x % 32;
    if (warp >= n_ids) return;
    // 64-bit offsets: 100M rows x 256 floats overflows 32 bits (and long is 32-bit on Windows)
    const float* src = host_table + ids[warp] * (int64_t)dim;
    float* dst = out + (int64_t)warp * dim;
    for (int j = lane; j < dim; j += 32) dst[j] = src[j];
}

Two details make this work. The output goes to HBM, so the rest of the training step reads it at full speed. And the sampled IDs should be sorted, which groups requests for neighbouring rows and lets the host's memory controller work on open DRAM pages.

Pattern 2: results the host reads while the GPU runs

The second pattern runs the other way: the GPU writes something small that the host must see soon, such as a progress counter, an early-exit flag or a handful of statistics. A zero-copy buffer avoids launching a copy for a few bytes. The writes are only guaranteed visible to the host after the kernel completes and you synchronise; if the host polls while the kernel runs, the kernel must order its writes with __threadfence_system() and the host must read through a volatile pointer so the compiler does not cache the value.

struct Mailbox { volatile int done; volatile float loss; };   // in mapped memory

__global__ void step(Mailbox* mb /*, ... */) {
    // ... compute ...
    if (blockIdx.x == 0 && threadIdx.x == 0) {
        mb->loss = 0.123f;
        __threadfence_system();   // make the payload visible before the flag
        mb->done = 1;
    }
}
// Host: while (!mb->done) { /* spin or sleep */ }  then read mb->loss

Atomics need more care. The CUDA programming guide warns that atomic operations on mapped page-locked memory are not atomic from the host's point of view; system-scope variants such as atomicAdd_system exist, but whether the host sees them as atomic depends on platform support you must query, so prefer one writer per location.

Write-combined buffers

Adding cudaHostAllocWriteCombined allocates the pages as write-combined: uncached on the CPU, with CPU writes merged into bursts. The runtime documentation says this can transfer across PCIe more quickly on some systems, and that CPU reads from such memory are very slow. Use it only for buffers the CPU writes sequentially and never reads, such as an input ring a producer thread fills for kernels to consume through zero-copy. Never use it for results the host reads.

Integrated and coherent systems

On integrated GPUs, such as the Jetson family, the CPU and GPU share the same DRAM, so copying between host and device buffers moves bytes from one part of the same memory to another. Zero-copy is often the right default there, though cache behaviour of pinned memory differs between Jetson generations; check NVIDIA's Jetson memory management guidance for your device. cudaDevAttrIntegrated tells you which case you are in.

Coherent CPU-GPU systems change the arithmetic. On Grace Hopper and Grace Blackwell, NVLink-C2C connects CPU and GPU with hardware coherence at 900 GB/s of total bandwidth according to NVIDIA, and the GPU can access CPU memory, including memory from plain malloc, at cache-line granularity. L is then much closer to H, so the break-even moves toward zero-copy for more patterns, but HBM is still several times faster and keeps the edge for data read repeatedly. The Grace Blackwell article covers the platform.

Failure modes

• A loop re-reads mapped memory. A stencil or reduction that touches each input several times runs at link speed. Nsight Compute's memory chart shows the traffic as system memory rather than device memory; Nsight Systems shows no copy at all, which is why this hides.
• Uncoalesced access. Scattered small reads waste the link. Restructure so warps read contiguous runs, or gather on the GPU with sorted IDs.
• Host reads before synchronisation. Results look correct in tests and stale under load. Synchronise the stream, or use the fence-and-flag pattern.
• Freeing or unregistering while a kernel runs. The GPU faults or reads garbage. Tie buffer lifetime to stream completion.
• Remote NUMA node. A buffer pinned on the other socket adds a hop to every access. Allocate from a thread bound to the GPU's local node.
• Over-pinning. Mapping a 100 GB table pins 100 GB the OS can no longer reclaim; size host memory and container limits for it.

What to do next

1. For each host-to-device transfer, estimate how many times the kernel reads each byte; anything well above one should stay a copy.
2. Look for sparse gathers from host-resident tables and CPU gather loops that precede a copy; those are the zero-copy candidates.
3. Allocate with cudaHostAllocMapped or register with cudaHostRegisterMapped once at start-up, and always use cudaHostGetDevicePointer.
4. Make every warp read contiguous 128-byte runs and sort gather IDs.
5. Profile with Nsight Compute and confirm system-memory throughput near link speed.
6. Synchronise or fence before the host reads anything a kernel wrote, and re-run the break-even on coherent or integrated systems instead of reusing PCIe conclusions.