Unified Memory

On Pascal+ (Linux), pages migrate on page faults; GPU memory can be oversubscribed (data larger than GPU memory).

19.1 Managed memory

On Pascal+ (Linux), pages migrate on page faults; GPU memory can be oversubscribed (data larger than GPU memory). On systems with HMM or Grace Hopper (coherent NVLink-C2C), even malloc-ed memory is GPU-accessible.

float* data;
cudaMallocManaged(&data, n * sizeof(float));      // accessible from CPU and GPU with one pointer
for (int i = 0; i < n; ++i) data[i] = i;          // CPU writes
scale<<<blocks, threads>>>(data, n, 2.0f);        // GPU reads/writes (pages migrate on demand)
cudaDeviceSynchronize();                           // REQUIRED before the CPU touches the data again
printf("%f\n", data[10]);
cudaFree(data);

__managed__ int counter;                           // statically allocated managed variable

19.2 Performance tuning: prefetch & advise

(CUDA 12.2+ adds cudaMemLocation-based variants of these APIs.)

int device; cudaGetDevice(&device);
cudaMemPrefetchAsync(data, bytes, device, stream);          // move pages to the GPU before the kernel
cudaMemPrefetchAsync(data, bytes, cudaCpuDeviceId, stream); // back to the CPU
cudaMemAdvise(data, bytes, cudaMemAdviseSetReadMostly, device);        // duplicate read-only pages
cudaMemAdvise(data, bytes, cudaMemAdviseSetPreferredLocation, device);
cudaMemAdvise(data, bytes, cudaMemAdviseSetAccessedBy, device);        // map without migrating

19.3 When to use unified memory

✅ Prototyping, complex pointer-based data structures (linked lists, trees), oversubscription, CPU/GPU-shared data. ⚠️ Page-fault overhead can hurt performance — use prefetching, or explicit memory management for peak performance.