CUDA Unified Memory is the programming model that provides a single, coherent address space accessible from both CPU and GPU — automatically migrating pages between CPU and GPU memory on demand, eliminating the need for explicit cudaMemcpy calls, simplifying GPU programming at the cost of potential performance overhead from page faults and migration latency.

Traditional vs. Unified Memory

How Unified Memory Works (Pascal+)

1. cudaMallocManaged(&ptr, size) — allocates in unified virtual address space. 2. Pages initially reside on CPU. 3. GPU kernel accesses ptr → page fault → GPU driver migrates page from CPU to GPU. 4. CPU accesses ptr → page fault → driver migrates page from GPU to CPU. 5. Pages migrated on demand at granularity of 4 KB (CPU page) or 64 KB (GPU preferred).

Performance Considerations

• First-touch penalty: Initial page fault and migration can be expensive (~20-50 μs per fault).
• Thrashing: If CPU and GPU both access same pages repeatedly → constant migration → terrible performance.
• Prefetching: cudaMemPrefetchAsync(ptr, size, device) — proactively migrate pages → avoids faults.
• Memory advise: cudaMemAdvise(ptr, size, advice, device) — hint system about access patterns.
• cudaMemAdviseSetReadMostly: Duplicate page on both CPU and GPU → no migration needed for reads.
• cudaMemAdviseSetPreferredLocation: Suggest where pages should reside.
• cudaMemAdviseSetAccessedBy: Allow remote access without migration.

Oversubscription

• Unified Memory enables GPU memory oversubscription — total allocation > GPU DRAM.
• Pages automatically evicted from GPU when GPU memory is full.
• Enables running workloads that don't fit in GPU memory (with performance penalty).
• Useful for: Prototyping, occasional large data, graceful degradation.

Performance Optimization Pattern

// Allocate managed memory
cudaMallocManaged(&data, N * sizeof(float));

// Initialize on CPU
initialize_data(data, N);

// Prefetch to GPU before kernel (avoids page faults during kernel)
cudaMemPrefetchAsync(data, N * sizeof(float), gpuDevice);

// Run kernel — data already on GPU, no faults
kernel<<<blocks, threads>>>(data, N);

// Prefetch back to CPU before CPU access
cudaMemPrefetchAsync(data, N * sizeof(float), cudaCpuDeviceId);
use_results(data, N);

When to Use Unified Memory

CUDA Unified Memory is an essential productivity tool that democratizes GPU programming — by removing the most error-prone aspect of GPU development (manual memory management), it enables faster development and handles complex data structures that would be impractical with explicit copies, while prefetching and memory advise hints allow recovering most of the performance.