Pinned (Page-Locked) Memory is host memory that is locked in physical RAM and cannot be swapped to disk — enabling the GPU to access host memory directly via DMA without CPU involvement and allowing asynchronous (overlapping) memory transfers.

Why Pinned Memory?

• Regular (pageable) memory: CPU can swap pages to disk. DMA transfer requires:

1. Allocate temporary pinned buffer. 2. Copy from pageable → pinned (CPU). 3. DMA transfer pinned → GPU.

• Double copy, synchronous.

• Pinned memory: Skip step 1-2 → DMA directly from host.
• 1.5–2x faster transfer bandwidth.
• Enables cudaMemcpyAsync — true asynchronous transfer.

Allocating Pinned Memory

float* h_data;
cudaMallocHost(&h_data, size);     // Pinned allocation
cudaFreeHost(h_data);              // Free pinned memory

// Async transfer (non-blocking)
cudaMemcpyAsync(d_data, h_data, size, cudaMemcpyHostToDevice, stream);

Zero-Copy Memory

• Map pinned host memory into GPU address space.
• GPU accesses host memory directly via PCIe (no explicit transfer).
• cudaHostAlloc(ptr, size, cudaHostAllocMapped)
• Useful when: Data accessed once (transfer + use = same latency as zero-copy), or host memory larger than GPU memory.
• Slower than transfer + compute: PCIe bandwidth ~16 GB/s vs. GPU memory ~900 GB/s.

When to Use Pinned Memory

• Always: For streaming/pipelined workloads with cudaMemcpyAsync.
• Large transfers: Bandwidth gain justifies pinning overhead.
• High-frequency small transfers: Saves per-transfer staging cost.

When NOT to Overuse

• Pinned memory cannot be swapped → reduces available virtual memory.
• Over-allocation: System runs low on physical memory → performance degradation.
• Rule: Pin only the buffers actively used for DMA transfers.

Pinned memory is a prerequisite for achieving peak PCIe bandwidth and enabling the transfer-compute overlap that allows GPU inference and training pipelines to saturate GPU compute without waiting for data transfers.