Pinned Memory & Data Transfer

malloc memory is pageable: the driver copies it into a pinned staging buffer before DMA.

18.1 Pageable vs pinned (page-locked) memory

malloc memory is pageable: the driver copies it into a pinned staging buffer before DMA. Allocating pinned memory directly gives higher bandwidth and enables async copies.

Don't over-allocate pinned memory — it reduces memory available to the OS.

float* h_pinned;
cudaMallocHost(&h_pinned, bytes);                       // or cudaHostAlloc(&p, bytes, flags)
cudaHostRegister(existing_ptr, bytes, cudaHostRegisterDefault);   // pin existing memory
cudaFreeHost(h_pinned);

18.2 Mapped (zero-copy) memory

Useful for data read once, integrated GPUs (Jetson), or when data doesn't fit on the GPU.

cudaHostAlloc(&h_ptr, bytes, cudaHostAllocMapped);
cudaHostGetDevicePointer(&d_ptr, h_ptr, 0);             // GPU accesses host memory over PCIe
kernel<<<...>>>(d_ptr);

18.3 Transfer optimization rules

Minimize host↔device transfers — keep data resident on the GPU across kernels.

Batch small transfers into large ones.

Use pinned memory and async copies to overlap with computation.

Use compression or lower precision for transfers when acceptable.

PCIe Gen4 x16 ≈ 25 GB/s, Gen5 ≈ 50 GB/s, NVLink-C2C (Grace Hopper) ≈ 900 GB/s — far below GPU memory bandwidth.

GPUDirect Storage / RDMA to bypass the CPU for storage and network I/O.

# Measure bandwidth: nvbandwidth (github.com/NVIDIA/nvbandwidth) or the bandwidthTest sample