Pinned and Constant Memory¶

Ordinary host memory allocated with malloc may be paged out by the operating system, so every copy from it to the device first goes through a pinned staging buffer. This episode introduces pinned (page-locked) host memory, which avoids that extra hop, can be accessed directly from device code and is required for asynchronous copies — together with the caveats that explain why not everything should be pinned.

The second part covers constant memory: a small (64 KB per device) read-only memory space with a dedicated cache on every streaming multiprocessor, meant for data that all threads read and nobody writes. A discussion at the end asks you to decide which of two arrays belongs in which of the two memory types.

Questions

Why is a copy out of ordinary host memory slower than the hardware would allow?
Pinning memory makes transfers faster — so why not pin everything?
Where should a small table that every thread reads, and nobody writes, actually live?

Objectives

Know when and how to use pinned (page-locked) and constant memory

Instructor note

10 min teaching
0 min exercises

Page-locked (pinned) memory¶

Regular host memory allocated with malloc may be paged out by the OS, i.e. the memory might be moved in physical memory to a different location at any time while the pointers of the program remain valid. When cudaMemcpy copies data from pageable memory, it must first copy it to a pinned staging buffer, meaning a location that is page-locked and cannot be moved to different physical location in memory, then transfer it over PCIe — this adds overhead.

Pinned (page-locked) memory is guaranteed to stay in physical memory and can be transferred directly, providing higher bandwidth.

The two diagrams below contrast the two transfer paths. Click the step buttons to compare.

Data transfer via staging buffer (pageable host memory) — Step 1 — Pageable host memory: CUDA copies into a hidden pinned staging buffer first, then transfers over PCIe.

Direct transfer with pinned host memory — Step 2 — Pinned host memory: transfer goes directly over PCIe, no staging hop.

Advantages:

Faster copy to/from device (no staging buffer needed)
Can be accessed directly from device code (with UVA, no explicit cudaHostGetDevicePointer needed)
Required for asynchronous memory copies (streams and asynchronous copies are covered in a later module)

Caveats

Be careful with race conditions when accessing mapped pinned memory from both host and device
CPU locks are not visible to the GPU; GPU atomics may not work on pinned host memory depending on the system
Access over PCIe is slow compared to device memory
CUDA Fortran does not allow pinned memory in kernel arguments and only allocatable arrays may be pinned

cudaError_t cudaHostAlloc(void** pHost, size_t size, unsigned int flags);
cudaError_t cudaMallocHost(void** ptr, size_t size);
cudaError_t cudaFreeHost(void* ptr);
// flags: cudaHostAllocMapped to map into device address space

real, allocatable, pinned :: q(:)
allocate(q(1024))

Note

Allocation/deallocation is analog to cudaMalloc/cudaFree (substitute with cudaMallocHost/cudaFreeHost) or device (replace with pinned).

Constant memory¶

Constant memory is a special read-only memory space that resides in global memory but uses a dedicated constant cache on each SM (8 KB). It is useful for data that is:

Read by all threads
Never written from device code
Small (maximum 64 KB per device)

Tip

Recall that variables passed by value to kernels are automatically stored in constant memory.

__constant__ float constData[256];

__global__ void myKernel() {
    float a = constData[0];  // Read from constant memory
}

// Host code: copy to/from constant memory
float data[256];
cudaMemcpyToSymbol(constData, data, sizeof(data), 0, cudaMemcpyDefault);
cudaMemcpyFromSymbol(data, constData, sizeof(data), 0, cudaMemcpyDefault);

module kernels
  real :: c_d(100)
  attributes(constant) :: c_d

  attributes(global) subroutine init()
    real :: c2
    c2 = c_d(2)  ! Read from constant memory
  end subroutine
end module

program main
  use kernels
  real :: c(100)
  c_d = c       ! Copy host → constant memory
end program

Restrictions

Cannot be dynamically allocated
Maximum 64 KB per device
Must not be written from device code
In C/C++ host code, constant memory variables are symbols — do not use pointer arithmetic on them

Which memory for which array?

Your kernel uses two arrays: a 100-entry table of coefficients that every thread reads and nobody writes, and a large input buffer copied to the device once per timestep.

One belongs in constant memory, the other is a candidate for pinned host memory. Which is which — and what goes wrong if you swap them?

Keypoints

Pinned memory enables faster transfers and is required for asynchronous copies
Constant memory provides cached, read-only access for small data (max 64 KB)