GPU I/O(3) — Unified Memory Architecture

🔍 Does CPU Memory Always Move into GPU HBM?

Not anymore. Conventional PCIe systems commonly migrate managed pages into GPU HBM before or during GPU execution. Grace Hopper and Grace Blackwell introduce another option: direct coherent access to CPU memory over NVLink-C2C.

---

① A Typical Unified Memory Application

From the application’s perspective, Unified Memory is simple. The application allocates managed memory, optionally provides memory hints, launches a kernel, and synchronizes when the work is complete.

float* data;

cudaMallocManaged(&data, size);
cudaMemAdvise(...);
cudaMemPrefetchAsync(...);

kernel<<<blocks, threads>>>(data);

cudaDeviceSynchronize();

Figure 1. Typical CUDA Unified Memory application flow. Although the application only invokes a few CUDA APIs, the runtime performs additional memory management behind the scenes.

At first glance, nothing here explains how the memory becomes accessible to the GPU.

That behavior depends on the underlying hardware.

---

② What Does cudaMemPrefetchAsync() Actually Do?

Many developers assume cudaMemPrefetchAsync() simply copies memory into GPU HBM. That is not what the API guarantees.

Instead, it tells CUDA Runtime:

“The GPU is expected to access this memory next. Prepare it for efficient GPU access.”

How that preparation happens depends on the platform.

Figure 2. Hardware behavior after cudaMemPrefetchAsync(). The application calls the same API on both systems, but CUDA Runtime prepares memory differently depending on the hardware.

---

③ PCIe vs Grace Superchip

On a conventional PCIe system, CPU DRAM and GPU HBM are separate memory domains.

If the managed page is not already resident in GPU memory, CUDA UVM commonly:

1. updates GPU page mappings,
2. migrates the page into GPU HBM,
3. resumes GPU execution.

The GPU typically accesses data only after the page becomes resident in HBM. Grace Hopper and Grace Blackwell keep exactly the same programming model. The difference is entirely in the hardware. Because Grace CPU and GPU are connected through coherent NVLink-C2C, the GPU can access CPU LPDDR directly without requiring page migration for correctness. The runtime may still migrate pages into HBM, but now migration is a performance optimization instead of a mandatory access path.

---

④ Why Does HBM Migration Still Exist?

Direct coherent access does not make HBM obsolete. HBM still provides much higher local bandwidth than remote CPU memory. CUDA Runtime therefore chooses between two access paths.

Remote coherent access

GPU
    ↓
NVLink-C2C
    ↓
Grace LPDDR

Local HBM access

GPU
    ↓
GPU HBM

Frequently reused GPU data benefits from HBM. Occasionally accessed or CPU-shared data may remain in Grace LPDDR. The application does not make this decision explicitly. CUDA Runtime chooses the most appropriate placement according to runtime behavior and programmer hints such as cudaMemPrefetchAsync() and cudaMemAdvise().

---

Key Takeaway

The Unified Memory API has not changed. The hardware implementation has.

PCIe GPU	Grace Superchip
Managed pages commonly migrate into GPU HBM	GPU can directly access coherent CPU LPDDR
Migration is commonly required before GPU execution	Migration becomes optional
cudaMemPrefetchAsync() often results in HBM migration	cudaMemPrefetchAsync() prepares the best placement for the platform

The same CUDA application runs on both systems.

Only the runtime’s memory management strategy changes.

---

TL;DR

• Unified Memory provides a unified virtual address space, not a single physical memory.
• cudaMemPrefetchAsync() is a placement hint rather than a copy command.
• PCIe systems commonly migrate managed pages into GPU HBM.
• Grace Superchip allows coherent GPU access to CPU LPDDR through NVLink-C2C.
• HBM migration still exists on Grace, but it is used for performance rather than correctness.