NVIDIA GH200, in depth: the Grace Hopper superchip, NVLink-C2C coherence and how software uses a CPU memory tier

What is on the module

For comparison, a PCIe Gen5 x16 link carries about 64 GB/s each way, so NVLink-C2C has roughly seven times its bandwidth. The two-superchip GH200 NVL2 node joins two modules with NVLink and is quoted by NVIDIA at 288 GB of HBM and 10 TB/s in total, consistent with the per-module HBM3e figures above. Because power is a range, two GH200 systems can perform quite differently; ask what limit a cloud instance runs at.

One address space over NVLink-C2C

In a normal server, the GPU has its own memory and its own address translation, and the CPU's memory is reachable only through explicit copies or slow page-faulting over PCIe. On GH200 the GPU uses Address Translation Services (ATS) to translate addresses through the same system page table the CPU uses. A pointer from malloc is therefore valid inside a CUDA kernel, and loads and stores to it travel over NVLink-C2C as coherent cache-line accesses, with no page fault and no copy.

CUDA's programming guide describes this as full unified memory support with hardware coherence. Three device attributes report it, and a program should check them rather than assume them:

int cma = 0, pageable = 0, hostPT = 0;
cudaDeviceGetAttribute(&cma,      cudaDevAttrConcurrentManagedAccess, dev);
cudaDeviceGetAttribute(&pageable, cudaDevAttrPageableMemoryAccess, dev);
cudaDeviceGetAttribute(&hostPT,   cudaDevAttrPageableMemoryAccessUsesHostPageTables, dev);
// 1, 1, 1  -> all system memory is unified, with hardware coherence (ATS): GH200
// 1, 1, 0  -> unified through software (HMM) on a PCIe system: works, but by page faults

Two boundaries remain. Memory from cudaMalloc is still GPU-only; the guide states that hardware coherence does not give the host access to it. Memory from cudaMallocManaged that resides on the GPU can be read by the CPU without migration.

Where pages actually live

Being reachable is not the same as being local, and performance depends on where pages physically live. Four mechanisms decide that.

• First touch. malloc only reserves addresses. A page gets physical memory when it is first written, on the side that writes it. A buffer initialised by a CPU loop lands in LPDDR5X, even if only the GPU will use it afterwards.
• Access-counter migration. The driver can count GPU accesses to CPU-resident pages and migrate hot ones into HBM. A 2024 study of Grace Hopper (arXiv 2407.07850) measured this with a default threshold of 256 accesses; treat that as driver behaviour that can change, not a specification.
• Page size. The OS page size, 4 KB or 64 KB on these systems, changes the cost of mapping and migration. The same study found 64 KB pages greatly reduced allocation overhead but could slow short kernels because each migration moves more data.
• Explicit placement. For data you know is hot, allocate it in HBM, through cudaMalloc or managed memory with prefetching and cudaMemAdvise hints, instead of hoping migration discovers it. Depending on driver configuration, the HBM may also appear to Linux as a separate NUMA node without CPUs; check with numactl --hardware.

The same study measured roughly 3.4 TB/s from HBM3, about 486 GB/s from LPDDR5X, and about 375 GB/s host-to-device and 297 GB/s device-to-host over the link, against 450 GB/s nominal. The practical rule: a byte read from the CPU tier costs the GPU about ten times as much as a byte from HBM. Use the CPU tier for data that is large and touched rarely, or read once per step, never for the inner loop. The HBM article explains why the HBM side is so fast.

A program that uses malloc memory on the GPU

A minimal program shows the model. The buffer is ordinary malloc memory, the kernel uses it directly, and the CPU reads the result without a copy.

__global__ void scale(float* x, size_t n, float a) {
    size_t i = blockIdx.x * (size_t)blockDim.x + threadIdx.x;
    if (i < n) x[i] *= a;
}

int main() {
    size_t n = 1ull << 30;                           // 4 GiB of floats
    float* x = (float*)malloc(n * sizeof(float));    // plain system allocation
    for (size_t i = 0; i < n; ++i) x[i] = 1.0f;      // CPU first touch: pages land in LPDDR5X
    scale<<<(n + 255) / 256, 256>>>(x, n, 2.0f);     // GPU reads and writes over NVLink-C2C
    cudaDeviceSynchronize();
    printf("%f\n", x[0]);                            // CPU reads the result directly
    free(x);
}

On an x86 server with a PCIe GPU and no HMM, this kernel would fault on an invalid address. On GH200 it is correct, and its speed depends on the C2C link because the pages start on the CPU side. If the GPU initialised the buffer instead, the pages would land in HBM. This is why porting code to GH200 is often easy to make correct and harder to make fast.

Training: optimizer state in CPU memory

Training uses the CPU tier for state that is large but touched once per step: optimizer moments and master weights. Take a full fine-tune of an 8-billion-parameter model with Adam in mixed precision. Weights and gradients in BF16 take 16 GB each; FP32 master weights and the two Adam moments take 32 GB each. That is 128 GB before activations, which does not fit the 96 GB HBM3 version.

Keep the 96 GB of FP32 optimizer state in LPDDR5X and the BF16 weights, gradients and activations in HBM. Each step, the optimizer reads and writes the 96 GB of state once. At the measured link rates that is about 96 / 375 plus 96 / 297 seconds, roughly 0.6 s per step. The forward and backward pass costs about 6 times parameters times tokens: for a 64k-token step, 6 x 8e9 x 65,536 = 3.1e15 FLOPs, about 6 s at half of an assumed H100 SXM-class dense BF16 peak of about 1 PFLOP/s. The offload adds about 10 percent, and less with larger steps or with the update overlapped with the next forward pass. On a PCIe system the same traffic would take several seconds per step.

The alternative is to run the optimizer on the 72 Grace cores, which read LPDDR5X locally at about 500 GB/s. Then only gradients and updated BF16 weights cross the link, about 32 GB per step. Which is faster depends on how well your framework's CPU optimizer uses Arm vector units; measure both.

Worked example: a KV cache tier for a 70B model

The most common inference use is a second tier for the KV cache. Consider serving a 70-billion-parameter model with grouped-query attention: 80 layers, 8 KV heads, head dimension 128, FP16 cache. Each token stores 2 x 80 x 8 x 128 x 2 bytes = 327,680 bytes, 320 KiB.

1. On the 144 GB version, FP8 weights take about 70 GB. Leaving about 14 GB for activations and runtime buffers gives about 60 GB of KV cache, room for roughly 185,000 tokens in HBM.
2. A 32k-token shared prefix, such as a long system prompt plus a document, takes 32,768 x 320 KiB = 10.7 GB. HBM holds about five of these alongside live requests.
3. LPDDR5X can hold several hundred gigabytes more, around 35 such prefixes after leaving room for the OS and the server process.
4. Bringing one prefix back over the link takes 10.7 GB / 375 GB/s, about 29 ms. Recomputing it costs about 2 x 70e9 x 32,768 = 4.6e15 FLOPs, several seconds even at high FP8 utilisation.

So a prefix-cache hit from the CPU tier is about two orders of magnitude cheaper than a miss, and the tier multiplies the number of cached prefixes several times over. Inference engines expose this as host-memory KV offload; LMCache with vLLM is one example. The same reasoning applies to embedding tables and vector indexes that exceed HBM.

When to choose GH200

GH200 suits workloads limited by memory capacity per GPU rather than by GPU count: long-context inference, retrieval and recommender models with huge tables, graph workloads, fine-tuning that would otherwise need many GPUs only for memory. For training that needs tensor parallelism across many GPUs, an NVLink domain matters more; see the NVLink and NVSwitch article.

Operating GH200 nodes

Operating GH200 nodes is mostly ordinary GPU operations with three additions. First, watch two memory pools, not one: export LPDDR5X usage per process alongside GPU memory, and alert when the CPU tier approaches its limit, because an out-of-memory kill of an inference server that was using host memory for its KV cache looks like a GPU problem but is not. Second, record the module power limit with every benchmark result; the same job can differ noticeably between a 700 W and a 1000 W configuration, and cloud providers do not always state which they use. Third, if your configuration exposes HBM as a CPU-less NUMA node, bind the memory of CPU-side helpers such as data loaders and tokenizers to the Grace node (for example with numactl --membind) so their host allocations do not land in HBM; with 72 cores per GPU there is usually capacity to spare, so data preprocessing that needed a separate CPU fleet on x86 can often run on the node itself.

Keep container images per architecture and tag them clearly. A mixed x86 and Arm fleet with one image name is a frequent source of start-up failures that look like driver errors.

Failure modes

• Silent remote access. Correct code reads its hot data from LPDDR5X at a tenth of HBM speed. Profile with Nsight Systems and check where large buffers were first touched.
• Migration thrash. Data used alternately by CPU and GPU moves back and forth. Pin it with placement hints or restructure the phases.
• Missing aarch64 builds. Python wheels, containers and native libraries built only for x86 fail to install or fall back to slow paths. Use NVIDIA's Arm containers and check every dependency early.
• Assuming H100 server numbers. Power limits, one GPU per node and different host bandwidth change throughput. Benchmark on the actual instance.
• Memory accounting. Tools that report only GPU memory hide CPU-tier use, and the OS needs its share of LPDDR5X. Budget both pools explicitly.

What to do next

1. Run the three device attribute checks on your instance and record driver, CUDA version and power limit.
2. Measure HBM, LPDDR5X and link bandwidth yourself, for example with NVIDIA's nvbandwidth tool, before trusting any table, including this one.
3. List your largest data structures and classify each as hot (keep in HBM) or large and cold (candidate for the CPU tier).
4. Check first-touch placement: initialise GPU-hot buffers on the GPU or allocate them explicitly.
5. For fine-tuning, try optimizer state in CPU memory and compare a GPU-side and a CPU-side optimizer step.
6. For inference, enable host-memory KV offload and measure prefix hit latency against recompute.
7. Confirm every dependency has an aarch64 build before migrating a production service.
8. Continue with the GPU memory hierarchy to place on-chip data well too.