CUDA Unified Memory
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
| Aspect | Traditional (Explicit) | Unified Memory |
|---|---|---|
| Allocation | `cudaMalloc` (GPU) + `malloc` (CPU) | `cudaMallocManaged` (single pointer) |
| Data transfer | `cudaMemcpy(dst, src, size, direction)` | Automatic page migration |
| Pointer sharing | Separate CPU/GPU pointers | Same pointer on both |
| Programmer effort | High (manage all transfers) | Low (system handles migration) |
| Performance | Optimal (programmer controls transfers) | Good (may have page fault overhead) |
| Oversubscription | Error if GPU memory exceeded | Data spills to CPU 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
| Use Case | Recommendation |
|---|---|
| Complex data structures (linked lists, trees) | Unified (explicit copy impractical) |
| Prototyping / rapid development | Unified (simplicity) |
| Production HPC / ML | Explicit (maximum control and performance) |
| GPU memory oversubscription | Unified (only option) |
| Multi-GPU with peer access | Unified (simplifies multi-GPU) |
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.