CUDA NUMA API 编程实践

基于 Intel Xeon Platinum 8470Q (双路 NUMA) + RTX 5090 实际环境。本文聚焦 CUDA 中与 NUMA 相关的 API：pinned memory 分配策略、cudaMemAdvise、cudaMemPrefetchAsync，以及如何通过 CPU 亲和性绑定优化 H2D/D2H 带宽。

1. 背景：为什么 NUMA 影响 GPU 程序

1.1 NUMA 架构简述

在双路 CPU 系统中，每个 CPU socket 拥有自己的内存控制器和本地 DDR5 内存。访问本地内存快（~70 ns），跨 socket 访问远端内存慢（~140 ns）。这种架构称为 NUMA (Non-Uniform Memory Access)——”非统一”指的是访问延迟因地址而异：

Socket 0 (NUMA node 0)               Socket 1 (NUMA node 1)
  ├── CPU 0-51, 104-155                ├── CPU 52-103, 156-207
  ├── DDR5 385 GB (本地)               ├── DDR5 387 GB (本地)
  └── PCIe domain 7e/7f                └── PCIe domain 97/d7 (GPU + NVMe)
         ↓                                      ↓
    本地访问: 10 (distance)               GPU 在这里!
         ↓                                      ↓
    跨 socket 访问: 21 (distance) ←── UPI ──→ 跨 socket 访问: 21

1.2 GPU 受 NUMA 影响的原因

GPU 通过 PCIe 连接到特定 NUMA 节点（本例为 node 1）。当 CUDA 程序：

调用 cudaMallocHost 分配 pinned memory——内存从当前线程所在的 NUMA 节点分配
如果线程运行在远端 NUMA (node 0)，H2D/D2H 需经过 UPI（Ultra Path Interconnect，Intel 双路 CPU 间的点对点互连总线）跨 socket 传输
NUMA distance 为 21 vs 10，约 2.1× 延迟惩罚

关键洞察：cudaMallocHost 不关心 GPU 在哪个 NUMA——它只关心调用线程在哪个 CPU。如果线程在 node 0 上调用 cudaMallocHost，分配的 pinned memory 就在 node 0 的 DDR5 上，而 GPU 在 node 1——每次 DMA 都必须穿过 UPI。这就是为什么理解并控制 NUMA 亲和性至关重要的原因。

2. Pinned Memory 的 NUMA 亲和性

2.1 `cudaMallocHost` 的默认行为

// 内存从当前线程的 NUMA 节点分配
float *h_buf;
cudaMallocHost(&h_buf, size);  // NUMA 亲和性取决于线程位置

2.2 使用 `cudaHostAlloc` 指定 NUMA 策略

// 可移植——允许从任意 NUMA 分配（降低性能但灵活）
cudaHostAlloc(&h_buf, size, cudaHostAllocPortable);

// 写入合并——优化 H2D 但牺牲随机读性能
cudaHostAlloc(&h_buf, size, cudaHostAllocWriteCombined);

// 组合使用
cudaHostAlloc(&h_buf, size,
    cudaHostAllocPortable | cudaHostAllocWriteCombined);

Flag	效果	适用场景
默认	线程所在 NUMA 分配	线程已绑定到 GPU NUMA
`cudaHostAllocPortable`	跨 NUMA 可映射	多 GPU 位于不同 NUMA
`cudaHostAllocWriteCombined`	绕过 L1/L2 cache	纯 H2D 传输（不读回）

3. CPU 亲和性控制

3.1 为什么需要

GPU 位于 NUMA node 1 (CPUs 52-103, 156-207)。程序应绑定到这些核，确保 pinned memory 从 node 1 分配。

3.2 通过 `taskset` 绑定

# 绑定到 GPU 所在 NUMA 节点 (node 1)
taskset -c 52-103,156-207 ./my_gpu_program

# 验证
taskset -c 52-103,156-207 nvidia-smi

3.3 通过 `numactl` 绑定

# 安装
apt install numactl

# 绑定 NUMA node 和内存策略
numactl --cpunodebind=1 --membind=1 ./my_gpu_program

# 查看拓扑
numactl --hardware

3.4 程序内查询 NUMA

# 查看 GPU 所在 NUMA node
cat /sys/bus/pci/devices/0000:98:00.0/numa_node
# 输出: 1

# 各 NUMA 节点的 CPU 范围
cat /sys/devices/system/node/node1/cpulist
# 输出: 52-103,156-207

4. CUDA Managed Memory 的 NUMA 优化

4.1 `cudaMemAdvise`

在统一内存 (Managed Memory) 模式下，给驱动 NUMA 放置提示：

// 分配统一内存
float *data;
cudaMallocManaged(&data, size);

// 告知驱动：这块内存主要在 GPU 0 上被访问
cudaMemAdvise(data, size, cudaMemAdviseSetPreferredLocation, 0);

// 告知驱动：这块内存会被 CPU 读取（只读）
cudaMemAdvise(data, size, cudaMemAdviseSetReadMostly, cudaCpuDeviceId);

4.2 `cudaMemPrefetchAsync`

主动将数据迁移到目标设备：

// 预取到 GPU 0（消除首次访问的 page fault 延迟）
cudaMemPrefetchAsync(data, size, 0, stream);

// 预取回 CPU（GPU 处理完毕后）
cudaMemPrefetchAsync(data, size, cudaCpuDeviceId, stream);

4.3 完整示例

#include <cuda_runtime.h>
#include <stdio.h>

#define N (1024 * 1024)

int main() {
    float *data;
    cudaMallocManaged(&data, N * sizeof(float));

    // 告知驱动：主要在 GPU 0 被访问
    cudaMemAdvise(data, N * sizeof(float),
                  cudaMemAdviseSetPreferredLocation, 0);

    // 预取到 GPU，避免首次 kernel launch 的 page fault
    cudaMemPrefetchAsync(data, N * sizeof(float), 0, 0);
    cudaDeviceSynchronize();

    // GPU kernel 计算...

    // 预取回 CPU
    cudaMemPrefetchAsync(data, N * sizeof(float), cudaCpuDeviceId, 0);
    cudaDeviceSynchronize();

    // CPU 读取结果
    for (int i = 0; i < 10; i++)
        printf("data[%d] = %f\n", i, data[i]);

    cudaFree(data);
    return 0;
}

编译：

nvcc -o managed_mem managed_mem.cu

5. 设备属性查询

RTX 5090 在统一内存方面的能力：

cudaDeviceGetAttribute(&v, cudaDevAttrConcurrentManagedAccess, 0);
// = 1 → 支持 concurrent managed access (Hopper+)

cudaDeviceGetAttribute(&v, cudaDevAttrPageableMemoryAccess, 0);
// = 0 → 不支持 pageable memory access (仅数据中心 GPU 支持)

cudaDeviceGetAttribute(&v, cudaDevAttrDirectManagedMemAccessFromHost, 0);
// = 0 → 不支持从 Host 直接访问 managed memory

6. 最佳实践速查

场景	推荐做法
单 GPU + H2D 为主	`cudaHostAlloc(WriteCombined)` + `taskset` 绑核
单 GPU + D2H 频繁	默认 `cudaMallocHost` + 确保线程在 GPU NUMA
多 GPU 不同 NUMA	`cudaHostAlloc(Portable)` + `cudaSetDevice` 切换
Managed Memory	`cudaMemAdvise` + `cudaMemPrefetchAsync` 主动放置
仅需知道 NUMA 信息	`cat /sys/bus/pci/devices/<bdf>/numa_node`