CUDA C编程权威指南-第九章：多GPU编程

系列导航：导读 | 上一篇：第8章 GPU加速库和OpenACC | 下一篇：第10章 程序实现的注意事项

系列第 9 篇。之前都在单 GPU 上做文章；这一章扩展到多 GPU：单机多设备管理、在多卡上划分计算、点对点通信与跨设备同步，再延伸到 CUDA-aware MPI 和集群，从单卡到多卡、单机到多机串成一条线。

---

前言：从单 GPU 到多 GPU

到目前为止，书中大部分示例都只使用一块 GPU。但在实际的高性能计算与数据中心环境中，单机多卡、多机多卡已成为常态。多 GPU 编程能让应用突破单卡算力与显存上限：通过聚合多块 GPU 的算力与带宽处理更大规模的问题，或通过任务/数据并行缩短总执行时间。这一章会说明：CUDA 为多 GPU 编程提供了完整支持——从单进程内多设备管理、点对点（P2P）内存访问与传输、到跨 GPU 的流与事件同步，再到基于 CUDA-aware MPI 与 GPUDirect RDMA 的集群级扩展。掌握多设备下的「当前设备」语义、P2P 的启用与使用、跨设备同步的 cudaStreamWaitEvent，以及多节点下的 MPI+GPU 数据交换，是写出可扩展、高性能多 GPU 应用的关键；本章与第 2 章（性能模型）、第 6 章（流与事件）紧密衔接。

---

一、本章在全书中的位置与学习目标

1.1 为什么要学「多 GPU 编程」

前几章我们主要在单 GPU 上优化：编程模型、执行模型、内存层次、流与并发、指令级原语、以及用库与 OpenACC 获得加速。需要搞清楚的是指出：在实际的 HPC 与数据中心中，单机多卡、多机多卡已是常态；多 GPU 编程能突破单卡算力与显存上限，通过聚合多块 GPU 的算力与带宽处理更大规模问题，或通过任务/数据并行缩短总时间。多 GPU 的第一步是正确理解「当前设备」：cudaSetDevice 决定后续 cudaMalloc、核函数启动、流与事件所绑定的设备；点对点（P2P） 则允许设备间直接传输或直接访问对方内存，在支持 NVLink/PCIe P2P 的拓扑上显著降低延迟、提高带宽。跨设备同步依赖 cudaStreamWaitEvent（流可等待另一设备上的事件）；多节点则依赖 CUDA-aware MPI 与 GPUDirect RDMA。因此，理解多设备管理、P2P 的检查与启用、跨设备同步模式、多设备上的数据划分与计算-通信重叠，以及集群扩展时的进程-设备映射与亲和性，是写出可扩展、高性能多 GPU 应用的关键一步。

1.2 学完本章，你应该能回答

学习目标	检验方式
理解当前设备对分配、启动、流与事件的约束	能说明 cudaSetDevice 后 cudaMalloc/核函数/流/事件与哪块设备绑定；为何核函数必须发往当前设备的流
掌握多设备下流、事件与 cudaMemcpy 的行为表	能说出哪些操作可跨设备、哪些会失败（如 EventRecord 跨设备失败、StreamWaitEvent 可跨设备）
掌握 P2P 的检查与启用：cudaDeviceCanAccessPeer、cudaDeviceEnablePeerAccess	能写出双向 P2P 启用的代码；说明「当前设备」与「peer 设备」的关系
掌握 cudaMemcpyPeer / cudaMemcpyPeerAsync 及 UVA 下 cudaMemcpyDefault	能说明何时用 Peer、何时可用 Default；P2P 传输与「经主机中转」的带宽差异
掌握跨设备同步：cudaStreamWaitEvent(streamB, eventA, 0)	能写出「设备 B 的流等待设备 A 的事件」的完整流程
理解多设备计算划分：每设备独立内存与流、数据按设备划分、锁页主机内存	能写出两卡并行向量加等示例；说明粗粒度与细粒度划分的取舍
理解计算与通信重叠：多流 + 异步 P2P + 事件同步	能描述「传输流」与「计算流」如何用事件衔接，以隐藏设备间传输延迟
理解 CUDA-aware MPI 与 GPUDirect RDMA 的定位	能说明为何直接传设备指针可简化代码并利于 RDMA；集群扩展时进程-设备映射
了解 CPU 与 GPU 亲和性 对多 GPU 性能的影响	能说出将线程/进程绑定到靠近目标 GPU 的 CPU 可减少延迟、提高带宽

1.3 博客阅读导图（本章架构）

第 9 章  多 GPU 编程
│
├── 二、从一个 GPU 到多 GPU（书 9.1 节）★ 重难点
│   ├── 在多 GPU 上执行（设备枚举、cudaSetDevice、当前设备语义）
│   ├── 多设备下的流、事件与内存拷贝行为表 ★ 易错点
│   ├── 点对点通信（检查与启用、cudaMemcpyPeer、UVA 与 P2P）
│   └── 多 GPU 间同步（cudaStreamWaitEvent 跨设备）
│
├── 三、多 GPU 间细分计算（书 9.2 节）
│   ├── 在多设备上分配内存与资源
│   ├── 数据划分与负载均衡（粗/细粒度、强扩展与通信开销）
│   └── 完整示例：多设备并行向量加法
│
├── 四、计算与通信的重叠（书 9.3 节）
│   ├── 多 GPU 上重叠传输与计算
│   └── 事件驱动与流水线/双缓冲模式
│
├── 五、使用 CUDA-aware MPI 交换数据（书 9.4 节）
│   ├── 传统 MPI 与 CUDA-aware MPI 的对比
│   └── 设备指针直接参与 MPI 通信
│
├── 六、在 GPU 加速集群上扩展应用（书 9.5 节）
│   ├── 进程/线程与 GPU 的映射、数据与计算划分
│   ├── GPUDirect RDMA 与 NCCL 简介
│   └── 强扩展效率与通信开销
│
├── 七、CPU 与 GPU 亲和性（书 9.6 节）
│   └── 绑定线程/进程到靠近目标 GPU 的 CPU
│
└── 八、本章小结与重难点回顾

---

二、从一个 GPU 到多 GPU（书 9.1 节）

本节对应书中对多 GPU 编程的入门：如何在程序中管理多块 GPU、在每块 GPU 上执行核函数，以及多设备下的流、事件与内存拷贝行为。多 GPU 编程的第一步是正确枚举设备、设置当前设备，并理解「当前设备」对后续分配与启动的约束。

2.1 在多 GPU 上执行

在单 GPU 程序中，默认使用设备 0；所有 cudaMalloc、核函数启动、流与事件的创建都发生在该设备上。若要使用多块 GPU，必须显式选择当前设备，并在每次为某块 GPU 分配资源或发起计算前调用 cudaSetDevice(deviceId)。

设备枚举与属性查询

使用 cudaGetDeviceCount 获取系统中可用的 CUDA 设备数量，然后对每个设备调用 cudaGetDeviceProperties 查询属性。典型用法与 CUDA 编程指南一致，如下所示：

#include <stdio.h>

#define CHECK(call)                                                                 \
    do {                                                                           \
        cudaError_t err = call;                                                     \
        if (err != cudaSuccess) {                                                   \
            printf("Error: %s:%d, code: %d, reason: %s\n", __FILE__, __LINE__, err, \
                   cudaGetErrorString(err));                                        \
            exit(1);                                                                \
        }                                                                          \
    } while (0)

int main() {
    int deviceCount = 0;
    CHECK(cudaGetDeviceCount(&deviceCount));

    for (int device = 0; device < deviceCount; device++) {
        cudaDeviceProp deviceProp;
        CHECK(cudaGetDeviceProperties(&deviceProp, device));
        printf("Device %d: %s, Compute %d.%d, SM count %d\n",
               device, deviceProp.name, deviceProp.major, deviceProp.minor,
               deviceProp.multiProcessorCount);
    }
    return 0;
}

编译与运行：nvcc -o deviceQuery deviceQuery.cu && ./deviceQuery。在多卡机器上会列出每块 GPU 的名称、计算能力与多处理器数量。

设备选择与当前设备

cudaSetDevice(int deviceId) 将指定设备设为当前主机线程的当前设备。此后，直到再次调用 cudaSetDevice 为止：

• cudaMalloc 在该设备上分配内存；
• 核函数启动（如 kernel<<<grid, block>>>(...)）在该设备上执行；
• cudaStreamCreate、cudaEventCreate 创建的流和事件与该设备关联。

这一章会说明：若不调用 cudaSetDevice，当前设备默认为 0。下面这段代码演示了在两块 GPU 上分别分配内存并启动同一核函数（与书中及 CUDA 编程指南 3.4.1.2 一致）：

__global__ void MyKernel(float *data) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    data[i] = (float)i;  // 示例：写入线程索引
}

int main() {
    const size_t size = 1024 * sizeof(float);

    cudaSetDevice(0);
    float *d_p0 = NULL;
    cudaMalloc(&d_p0, size);
    MyKernel<<<4, 256>>>(d_p0);

    cudaSetDevice(1);
    float *d_p1 = NULL;
    cudaMalloc(&d_p1, size);
    MyKernel<<<4, 256>>>(d_p1);

    cudaSetDevice(0);
    cudaFree(d_p0);
    cudaSetDevice(1);
    cudaFree(d_p1);
    return 0;
}

要点：每次切换设备后，后续的分配与启动都作用于新当前设备；流和事件也与创建时所在设备绑定。

理解与体会：「当前设备」是主机线程的局部状态，不是全局的。同一进程内不同主机线程可以各自 cudaSetDevice(i) 控制不同的当前设备，从而实现多线程各自驱动一块 GPU；若单线程使用多 GPU，则必须在每次为某块 GPU 分配资源或发起计算前显式切换当前设备，否则极易出现「以为在 GPU 1 上分配，实际在 GPU 0 上」的隐蔽错误。

2.1.1 多设备下的流、事件与内存拷贝 ★ 易错点

CUDA 编程指南强调以下行为（多设备下流、事件与内存拷贝的规则，下表与之对应）：

操作	多设备下的行为
核函数启动	必须发往与当前设备关联的流；若流属于另一设备，启动会失败。
cudaMemcpy	即使当前设备与流所属设备不同，跨设备拷贝仍可成功（运行时会在正确的设备上执行拷贝）。
cudaEventRecord	若事件与流属于不同设备，则 失败。
cudaEventElapsedTime	若两个事件属于不同设备，则 失败。
cudaEventSynchronize / cudaEventQuery	即使事件属于非当前设备，也可成功（用于跨设备同步）。
cudaStreamWaitEvent	即使流与事件属于不同设备也可成功，因此可用于跨设备同步。

因此，cudaStreamWaitEvent 是实现「设备 B 上的流等待设备 A 上的事件」的关键 API，书中在多 GPU 同步一节会用到。

每块设备有各自的默认流；发往设备 0 的默认流的命令与发往设备 1 的默认流的命令可以乱序或并发执行，彼此不阻塞。

易错点：核函数启动必须发往与当前设备关联的流。若当前设备是 0 却把 kernel 发往在设备 1 上创建的流，启动会失败。cudaEventRecord 时若事件与流属于不同设备则失败；cudaEventElapsedTime 若两个事件属于不同设备则失败。跨设备同步只能依赖「流等待另一设备上的事件」——即 cudaStreamWaitEvent，这是多 GPU 同步的核心 API，务必熟记。

---

2.2 点对点通信

在多 GPU 之间交换数据时，若不经特殊处理，数据往往需要先从一个 GPU 拷回主机，再从主机拷到另一块 GPU（两次 PCIe 传输、经过主机内存）。点对点（Peer-to-Peer, P2P） 机制允许两块 GPU 直接在设备内存之间传输数据（或直接访问对方内存），在支持 P2P 的拓扑（如通过 NVLink 或同一 PCIe 根下的 PCIe P2P）上可以显著降低延迟并提高带宽。书中将点对点分为点对点传输和点对点内存访问两类。

检查与启用点对点访问

cudaDeviceCanAccessPeer(int *canAccessPeer, int device, int peerDevice)：查询 device 是否可以直接访问 peerDevice 的设备内存。若拓扑支持（例如两块 GPU 通过 NVLink 或 PCIe P2P 相连），则 *canAccessPeer == 1。

cudaDeviceEnablePeerAccess(int peerDevice, unsigned int flags)：在当前设备上启用对 peerDevice 的单向点对点访问。调用前需先 cudaSetDevice(currentDevice)；flags 通常传 0。这一章会说明：启用后，当前设备上的核函数可以直接读写 peerDevice 上已分配的内存（需在支持 P2P 访问的架构上）。若需双向访问，两个设备上都要调用一次（例如设备 0 对设备 1 启用、设备 1 对设备 0 启用）。

典型用法（与书中及编程指南一致）：

int canAccessPeer01 = 0, canAccessPeer10 = 0;
cudaDeviceCanAccessPeer(&canAccessPeer01, 0, 1);
cudaDeviceCanAccessPeer(&canAccessPeer10, 1, 0);

if (canAccessPeer01 && canAccessPeer10) {
    cudaSetDevice(0);
    cudaDeviceEnablePeerAccess(1, 0);
    cudaSetDevice(1);
    cudaDeviceEnablePeerAccess(0, 0);
}

点对点传输：cudaMemcpyPeer

设备间拷贝可以使用 cudaMemcpyPeer（同步）或 cudaMemcpyPeerAsync（异步，需指定流）：

cudaError_t cudaMemcpyPeer(void *dst, int dstDevice, const void *src, int srcDevice, size_t count);
cudaError_t cudaMemcpyPeerAsync(void *dst, int dstDevice, const void *src, int srcDevice,
                                size_t count, cudaStream_t stream);

当 P2P 访问已启用 时，这类拷贝会走设备间直连路径（如 NVLink 或 PCIe P2P），而不经过主机，带宽与延迟明显优于「Device→Host→Device」。在支持 UVA（统一虚拟寻址）的平台上，若已启用 P2P，也可以使用 cudaMemcpy(..., cudaMemcpyDeviceToDevice) 或 cudaMemcpy(..., cudaMemcpyDefault)，运行时可根据指针推断源与目标设备并选择 P2P 路径。

统一虚拟寻址（UVA）与 P2P

在 64 位且支持 UVA 的系统中，主机与所有设备的地址处在同一虚拟地址空间。因此，通过指针值即可判断某地址属于哪块设备（或主机）。启用 P2P 后，从一块 GPU 上运行的核函数中，可以直接解引用指向另一块 GPU 内存的指针进行 load/store（即点对点内存访问），无需显式拷贝。此时同一指针可在不同设备上使用（例如在设备 1 上启用对设备 0 的 P2P 后，设备 1 的核函数可以接收并访问设备 0 上分配的指针）。UVA 下使用 cudaMemcpy 做设备间拷贝时，若使用 cudaMemcpyDefault，源和目标会由运行时根据地址自动识别。

方式	API	说明
显式指定设备	cudaMemcpyPeer(dst, dstDev, src, srcDev, size)	不依赖 UVA，适用于所有支持 P2P 传输的环境
UVA + P2P	cudaMemcpy(dst, src, size, cudaMemcpyDefault)	需 UVA；运行时根据指针推断设备，若已启用 P2P 则走直连

理解与体会：P2P 分为传输与访问两种用法。传输用 cudaMemcpyPeer 或 UVA 下的 cudaMemcpy，数据从一卡拷到另一卡而不经主机，带宽与延迟优于「Device→Host→Device」。访问则是在启用 P2P 后，某设备上的核函数直接解引用指向对端设备内存的指针做 load/store，适合细粒度、不规则访问模式；此时需确保对端内存在核函数执行期间有效且语义正确（避免竞态）。多数场景下先掌握 P2P 传输即可，P2P 访问在跨设备共享只读表或边界数据时很有用。

---

2.3 多 GPU 间同步

多 GPU 应用中，经常需要「等某块 GPU 上的计算或传输完成后再在另一块 GPU 上继续」。由于流和事件与设备绑定，同步应通过事件 + 流等待完成，而不是假设默认流的隐式顺序。

典型模式如下：

1. 在设备 A 上创建流 streamA 和事件 eventA；
2. 在设备 A 的 streamA 中执行核函数或拷贝，然后 cudaEventRecord(eventA, streamA)；
3. 在设备 B 上创建流 streamB；
4. 在设备 B 上执行 cudaStreamWaitEvent(streamB, eventA, 0)，使 streamB 在 eventA 完成之后再执行后续操作；
5. 在 streamB 中启动核函数或拷贝。

这样，设备 B 上的工作会等待设备 A 上 eventA 之前的工作完成。这一章会说明：cudaStreamWaitEvent 即使流与事件属于不同设备也能成功，因而是跨 GPU 同步的核心 API。

下面是一段跨设备同步的完整示例（风格与书中一致）：设备 0 上先执行核函数并记录事件，设备 1 上的流等待该事件后再执行核函数。

__global__ void kernelA(float *out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) out[i] = 1.0f;
}

__global__ void kernelB(float *in, float *out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) out[i] = in[i] * 2.0f;
}

int main() {
    const int n = 1024;
    size_t bytes = n * sizeof(float);

    cudaSetDevice(0);
    float *d_A0 = NULL, *d_B0 = NULL;
    cudaMalloc(&d_A0, bytes);
    cudaMalloc(&d_B0, bytes);
    cudaStream_t s0;
    cudaEvent_t e0;
    cudaStreamCreate(&s0);
    cudaEventCreate(&e0);

    cudaSetDevice(1);
    float *d_B1 = NULL;
    cudaMalloc(&d_B1, bytes);
    cudaStream_t s1;
    cudaStreamCreate(&s1);

    // 设备 0：在 s0 中执行 kernelA，然后记录 e0
    cudaSetDevice(0);
    kernelA<<<(n + 255) / 256, 256, 0, s0>>>(d_B0, n);
    cudaEventRecord(e0, s0);

    // 设备 1：s1 等待 e0 后再执行 kernelB
    cudaSetDevice(1);
    cudaStreamWaitEvent(s1, e0, 0);
    kernelB<<<(n + 255) / 256, 256, 0, s1>>>(d_B1, d_B1, n);  // 此处 d_B1 仅作示例

    cudaDeviceSynchronize();  // 等待所有设备
    cudaSetDevice(0);
    cudaStreamDestroy(s0);
    cudaEventDestroy(e0);
    cudaFree(d_A0);
    cudaFree(d_B0);
    cudaSetDevice(1);
    cudaStreamDestroy(s1);
    cudaFree(d_B1);
    return 0;
}

---

三、多 GPU 间细分计算（书 9.2 节）

本节对应书中在多块 GPU 上分配工作与数据的方法：如何在多设备上分配内存、如何将问题划分到各 GPU，以及如何结合 P2P 或主机中转进行数据交换。书中区分粗粒度（每块 GPU 独立处理一块子问题）和细粒度（更细的数据划分与负载均衡）两种思路；粗粒度实现简单，细粒度有利于负载均衡但通信与协调成本更高。

3.1 在多设备上分配内存

多 GPU 程序通常需要为每块设备单独分配：

• 设备内存：cudaSetDevice(i) 后 cudaMalloc(&d_ptr_i, size)；
• 主机端资源：若使用异步传输，常用锁页主机内存（cudaMallocHost）以支持与多设备的异步拷贝；
• 流和事件：每个设备创建自己的流和事件，用于该设备上的并发与同步。

典型模式：用循环对每个设备执行「设置当前设备 → 分配设备内存、创建流/事件」；主机端可分配一块或每设备一块锁页缓冲区，用于与各设备的 H2D/D2H 传输。数据划分时，常见做法是将总数据量按设备数均分（或按负载均衡策略分配），每块 GPU 处理自己的那一份；若有依赖（例如归约），则再通过 P2P 或主机做一次汇总。

下表归纳多设备下的资源与 API 使用：

资源类型	多设备下的做法
设备内存	每设备 cudaSetDevice(i) 后 cudaMalloc，指针分别保存（如 d_A[0], d_A[1]）
主机内存	锁页 cudaMallocHost 便于与多设备异步拷贝；或每设备一块 sub-buffer
流 / 事件	每设备创建自己的流和事件，核函数与拷贝发往对应设备的流
数据划分	按设备数均分（或按负载分配），每设备处理自己的区间

两卡环境下按数据均分的逻辑可概括为（每块 GPU 处理一段连续子数据，结果写回设备内存，最后由主机合并或再经 P2P 归约）：

flowchart LR
    subgraph Host [主机]
        H_Mem[(锁页内存\n完整数组)]
    end
    subgraph GPU0 [GPU 0]
        D0[(设备内存\n前半段)]
        K0[核函数]
    end
    subgraph GPU1 [GPU 1]
        D1[(设备内存\n后半段)]
        K1[核函数]
    end
    H_Mem -->|H2D 前半| D0
    H_Mem -->|H2D 后半| D1
    D0 --> K0
    D1 --> K1
    K0 --> D0
    K1 --> D1
    D0 -->|D2H| H_Mem
    D1 -->|D2H| H_Mem

3.2 完整示例：多设备上并行向量加法

下面这段代码演示在两块 GPU 上按数据划分做向量加法：主机分配锁页内存与设备内存，每个设备负责一半元素，最后在主机上合并结果（书中常见风格，对应多设备分配与执行）：

#include <stdio.h>
#include <stdlib.h>

#define CHECK(call)                                                                 \
    do {                                                                           \
        cudaError_t err = call;                                                     \
        if (err != cudaSuccess) {                                                   \
            printf("Error: %s:%d, reason: %s\n", __FILE__, __LINE__,                \
                   cudaGetErrorString(err));                                        \
            exit(1);                                                                \
        }                                                                          \
    } while (0)

__global__ void addKernel(const float *A, const float *B, float *C, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) C[i] = A[i] + B[i];
}

int main() {
    const int N = 1 << 20;           // 总元素数
    const int nBytes = N * sizeof(float);
    int nDevices = 0;
    CHECK(cudaGetDeviceCount(&nDevices));
    if (nDevices < 2) {
        printf("Need at least 2 GPUs.\n");
        return 0;
    }

    // 主机锁页内存（整块数据）
    float *h_A = NULL, *h_B = NULL, *h_C = NULL;
    CHECK(cudaMallocHost(&h_A, nBytes));
    CHECK(cudaMallocHost(&h_B, nBytes));
    CHECK(cudaMallocHost(&h_C, nBytes));
    for (int i = 0; i < N; i++) {
        h_A[i] = 1.0f;
        h_B[i] = 2.0f;
    }

    const int nHalf = N / 2;
    const size_t halfBytes = nHalf * sizeof(float);

    float *d_A[2], *d_B[2], *d_C[2];
    cudaStream_t st[2];

    for (int dev = 0; dev < 2; dev++) {
        CHECK(cudaSetDevice(dev));
        CHECK(cudaMalloc(&d_A[dev], halfBytes));
        CHECK(cudaMalloc(&d_B[dev], halfBytes));
        CHECK(cudaMalloc(&d_C[dev], halfBytes));
        CHECK(cudaStreamCreate(&st[dev]));
    }

    // 每设备：拷贝自己那一半，计算，拷回
    for (int dev = 0; dev < 2; dev++) {
        CHECK(cudaSetDevice(dev));
        int offset = dev * nHalf;
        CHECK(cudaMemcpyAsync(d_A[dev], h_A + offset, halfBytes, cudaMemcpyHostToDevice, st[dev]));
        CHECK(cudaMemcpyAsync(d_B[dev], h_B + offset, halfBytes, cudaMemcpyHostToDevice, st[dev]));
        addKernel<<<(nHalf + 255) / 256, 256, 0, st[dev]>>>(d_A[dev], d_B[dev], d_C[dev], nHalf);
        CHECK(cudaMemcpyAsync(h_C + offset, d_C[dev], halfBytes, cudaMemcpyDeviceToHost, st[dev]));
    }

    CHECK(cudaSetDevice(0));
    CHECK(cudaDeviceSynchronize());
    CHECK(cudaSetDevice(1));
    CHECK(cudaDeviceSynchronize());

    // 验证
    int ok = 1;
    for (int i = 0; i < N; i++) if (h_C[i] != 3.0f) { ok = 0; break; }
    printf("%s\n", ok ? "PASS" : "FAIL");

    for (int dev = 0; dev < 2; dev++) {
        cudaSetDevice(dev);
        cudaFree(d_A[dev]);
        cudaFree(d_B[dev]);
        cudaFree(d_C[dev]);
        cudaStreamDestroy(st[dev]);
    }
    cudaFreeHost(h_A);
    cudaFreeHost(h_B);
    cudaFreeHost(h_C);
    return 0;
}

编译：nvcc -o multiGPU_add multiGPU_add.cu。运行需至少两块 GPU。该示例体现了：按设备划分数据、每设备独立流、异步传输与计算、最后双设备同步的典型多 GPU 流程。

理解与体会：多设备编程的「标准模板」是：为每个设备单独 cudaSetDevice(i) → cudaMalloc、cudaStreamCreate（及必要时 cudaEventCreate），用数组或结构保存各设备的指针与流；主机端若需与多设备异步拷贝，使用锁页内存（cudaMallocHost）以便多流并发执行 H2D/D2H；数据按设备划分后，每设备只处理自己的区间，最后在主机合并或通过 P2P 做归约。先跑通「每设备独立流 + 异步拷贝 + 各设备同步」的流程，再考虑 P2P 与计算-通信重叠。

3.3 数据划分与负载均衡

这一章会说明：粗粒度划分（例如均匀分块）实现简单，但若各块工作量差异大，会出现负载不均衡。细粒度划分（例如按行/按任务动态分配）可以更好地平衡负载，但需要更多的设备间通信或主机协调。在强扩展（固定问题规模、增加 GPU 数）下，理想加速比可写为：

[
S_p = \frac{T_1}{T_p} \leq \frac{T_1}{T_1/p + T_{\text{comm}}} \approx \frac{p}{1 + p \cdot T_{\text{comm}}/T_1}
]

其中 (T_1) 为单 GPU 时间，(T_p) 为 (p) 块 GPU 的时间，(T_{\text{comm}}) 为通信与同步开销。当 (T_{\text{comm}}) 不可忽略时，加速比会低于 (p)；因此多 GPU 编程中减少设备间通信、重叠计算与通信是性能关键，与下一节「计算与通信的重叠」以及第 6 章流的思路一致。

---

四、计算与通信的重叠（书 9.3 节）

与单 GPU 上重叠主机-设备传输与核函数执行类似，在多 GPU 上应尽量重叠设备间传输与各设备上的计算。第 2 章曾提到：若计算时间大于数据传输时间，可通过流来重叠计算与通信、隐藏传输延迟。在多 GPU 场景下，这一思想同样适用：当 GPU 0 正在计算时，GPU 1 与 GPU 0 之间的 P2P 传输可以在另一流中进行；或当 GPU 0 把结果通过 P2P 发给 GPU 1 时，GPU 1 可以在收到一部分数据后就开始计算（双缓冲/流水线），从而提高整体吞吐。

实现方式与第 6 章一致：使用 异步拷贝（cudaMemcpyPeerAsync、cudaMemcpyAsync）和多个流，让传输与核函数分属不同流，由运行时在硬件允许的情况下并发执行。跨设备时，用 cudaStreamWaitEvent 让「计算流」等待「传输完成事件」，再启动依赖该数据的核函数。概括思路：每设备多流 + 事件同步 + 异步 P2P 是实现多 GPU 上计算-通信重叠的标准做法。

理解与体会：多 GPU 上的「重叠」与单 GPU 上「传输与计算重叠」是同一思想：用异步操作和事件把依赖关系表达清楚，让硬件在满足依赖的前提下尽量并行。设备 A 的「发送流」与设备 B 的「接收流」配合 P2P 传输；设备 B 上用 cudaStreamWaitEvent(computeStream, recvDoneEvent, 0) 再启动核函数，即可保证「算的是已收到的数据」。若有多阶段计算与通信（如迭代 stencil），可采用双缓冲：一块缓冲区在接收，另一块供当前核函数使用，下一轮交换角色，形成流水线。

下表归纳多 GPU 上「计算与通信重叠」的典型模式（与书中思想一致）：

模式	做法	目的
传输与计算并行	设备 A 的流 1 做 P2P 发送，流 2 做本地核函数；设备 B 的流 1 做 P2P 接收，流 2 做依赖数据的核函数	隐藏设备间传输延迟
流水线/双缓冲	当一块缓冲区在传输时，另一块用于计算；交替使用	提高设备利用率
事件驱动	用 cudaEventRecord 标记传输完成，cudaStreamWaitEvent 让计算流等待	保证数据就绪后再计算，避免竞态

---

五、使用 CUDA-aware MPI 交换数据（书 9.4 节）

在多节点集群中，GPU 分布在不同主机上，设备间无法直接 P2P，必须借助消息传递。CUDA-aware MPI 是指 MPI 实现能够直接使用 GPU 指针作为发送/接收缓冲区：在调用 MPI_Send/MPI_Recv（或 MPI_Isend/MPI_Irecv）时传入设备指针，MPI 库内部会与 CUDA 协作，在 GPU 内存与网络之间搬运数据，无需程序员先把数据拷回主机再发送。这一章会说明：这样既简化了代码，又便于与 GPUDirect RDMA 结合，减少主机内存与 CPU 参与，提高带宽并降低延迟。

典型用法（概念性）：各进程绑定一块 GPU（cudaSetDevice），在设备上分配发送/接收缓冲区，然后直接以设备指针调用 MPI 的发送/接收接口（具体 API 以所用 MPI 文档为准）。支持 CUDA-aware 的 MPI 实现包括 MVAPICH2、Open MPI、IBM Platform MPI 等；编译与链接时需启用对应选项（如 --with-cuda）并链接 CUDA 运行时。多节点矩阵-向量乘或 Jacobi 等示例中，在各 rank 上设置设备、分配设备缓冲区并调用 MPI 集体通信或点对点通信即可。

方式	说明
传统 MPI	进程先将 GPU 数据拷回主机，再用主机缓冲区调用 MPI_Send/Recv；接收后再拷到 GPU
CUDA-aware MPI	直接以设备指针调用 MPI；由 MPI 与 CUDA 协同完成与网络的传输

理解与体会：传统 MPI 下，若数据在 GPU 上，必须先 D2H 拷回主机、用主机缓冲区 MPI_Send/Recv、再 H2D 拷到目标进程的 GPU，不仅代码繁琐，而且多了一次主机参与和内存拷贝。CUDA-aware MPI 允许直接传设备指针，MPI 库内部与 CUDA 协作（在支持 GPUDirect RDMA 的环境中还可让网卡直接访问 GPU 内存），既简化了代码，又为减少拷贝、提高带宽创造条件。使用前需确认所用 MPI 实现支持并已启用 CUDA 支持（如 Open MPI 的 --with-cuda、MVAPICH2 的对应选项），编译链接时也需正确链接 CUDA 运行时。

---

六、在 GPU 加速集群上扩展应用（书 9.5 节）

书中将多 GPU 编程从单节点多卡推广到多节点 GPU 集群：每节点可有多块 GPU，节点间通过 MPI 通信。扩展应用时需要考虑：

• 进程/线程与 GPU 的映射：通常每进程或每线程绑定一块 GPU（cudaSetDevice），MPI 秩与设备号的对应关系可由用户或运行时设置。
• 数据与计算划分：按 MPI 进程（或按进程内多 GPU）划分域或任务，每进程负责本节点的 GPU 计算与本地通信，跨节点用 MPI 交换边界或全局数据。
• GPUDirect RDMA：在支持 InfiniBand 等网络的集群上，GPUDirect RDMA 允许网卡直接从 GPU 内存 DMA 到网络（或反向），无需经过主机内存。与 CUDA-aware MPI 结合后，多节点间的 GPU 到 GPU 传输可以接近线速，其优势与使用条件（驱动、硬件、MPI 实现支持）可参考原书或文档。

集群扩展的理想强扩展效率可表示为：

[
\text{Efficiency} = \frac{S_p}{p} = \frac{T_1}{p \cdot T_p} \leq 1
]

效率越接近 1，说明并行扩展越好；通信与负载不均衡会拉低效率。

这一章会说明：实际集群上还需考虑 NCCL（NVIDIA Collective Communications Library）等库，它们针对多 GPU 与多节点的集体操作（如 All-Reduce、Broadcast）做了优化，可与 MPI 配合或替代部分 MPI 集体调用，以获得更高的多 GPU/多节点通信性能。

理解与体会：集群扩展时，扩展效率 (E_p = S_p/p = T_1/(p \cdot T_p)) 越接近 1 越好；通信与负载不均衡会拉低效率。单节点内用 P2P 与多流重叠；跨节点则依赖 MPI（及可选 NCCL）与 GPUDirect RDMA。进程与 GPU 的映射（如每进程一 GPU、或 MPI rank 与 device id 的对应）需在程序初始化时明确设置（如通过 cudaSetDevice(rank % ngpus) 或环境变量 CUDA_VISIBLE_DEVICES 限制每进程可见的 GPU），避免多进程争抢同一块 GPU 或设备号与 NUMA 拓扑不匹配。

---

七、CPU 与 GPU 亲和性（书 9.6 节）

CPU 与 GPU 亲和性指的是：在多 GPU、多 NUMA 节点或多核主机上，将主机线程或进程绑定到靠近某块 GPU 的 CPU 核心，可以减少 PCIe 访问延迟、提高传输带宽，并避免操作系统将线程迁移到远离该 GPU 的 CPU。

在 Linux 上可通过 numactl、pthread_setaffinity_np 等设置 CPU 亲和性；部分 MPI 实现也支持绑定进程到指定 CPU 与 GPU。例如，在单节点多卡上，可以让「控制 GPU 0 的主机线程」绑定到与 GPU 0 同 NUMA 节点的 CPU，从而减少跨 NUMA 访问带来的延迟。具体 API 或环境变量（如 CUDA_VISIBLE_DEVICES 与进程绑定的配合）可参考原书或文档。理解亲和性有助于在多 GPU 与多节点环境中获得更稳定、更高的性能。

理解与体会：在多 NUMA 节点、多 PCIe 根的主机上，GPU 与 CPU 的「距离」并不相同：某块 GPU 可能挂在某颗 CPU 的 PCIe 下，若控制该 GPU 的主机线程被调度到另一颗 CPU，则 PCIe 访问可能经过跨节点互联，延迟与带宽都会受影响。通过将线程/进程绑定到与目标 GPU 同 NUMA 节点的 CPU，可减少这类问题。在实际部署时，可结合 nvidia-smi topo -m 查看 GPU 与 CPU 的拓扑，再配合 numactl 或 MPI 的绑定选项做亲和性设置。

---

八、本章小结与重难点回顾

本章系统梳理了 多 GPU 编程 从单节点多卡到集群扩展的核心内容，严格对应书中第 9 章各节。下面用表格与要点形式做小结，便于复习与自测。

8.1 知识点总览

主题	要点
多设备管理	cudaGetDeviceCount、cudaSetDevice；分配、流、事件与当前设备绑定；核函数必须发往当前设备的流。
点对点	cudaDeviceCanAccessPeer、cudaDeviceEnablePeerAccess；cudaMemcpyPeer / cudaMemcpyPeerAsync；UVA 下可用 cudaMemcpyDefault 做设备间拷贝。
跨设备同步	cudaStreamWaitEvent 可跨设备使用，实现「设备 B 的流等待设备 A 的某事件」；每设备有独立默认流。
多设备计算	每设备单独分配内存与流；数据按设备划分；锁页主机内存便于多设备异步传输。
计算与通信重叠	多流 + 异步 P2P 传输 + 事件同步，在各 GPU 上重叠传输与计算。
CUDA-aware MPI	直接使用设备指针进行 MPI 通信；集群扩展时结合 GPUDirect RDMA 提升跨节点 GPU 传输性能。
亲和性	将主机线程/进程绑定到靠近目标 GPU 的 CPU，可改善多 GPU 性能。

8.2 重难点速记

• 当前设备：cudaSetDevice(i) 后，后续的 cudaMalloc、核函数启动、cudaStreamCreate/cudaEventCreate 都作用于设备 i；核函数若发往非当前设备的流会失败。
• 多设备下易错：cudaEventRecord(event, stream) 若 event 与 stream 属不同设备则失败；cudaEventElapsedTime 若两事件属不同设备则失败。跨设备同步只能用 cudaStreamWaitEvent(streamB, eventA, 0)。
• P2P：先 cudaDeviceCanAccessPeer 检查，再在各设备上 cudaDeviceEnablePeerAccess(peer, 0) 启用（双向需两设备各启用一次）；传输用 cudaMemcpyPeer/cudaMemcpyPeerAsync，UVA 下可用 cudaMemcpy(..., cudaMemcpyDefault)。
• 多设备模板：每设备 cudaSetDevice → 分配设备内存与流（及事件）→ 数据按设备划分 → 异步 H2D → kernel → 异步 D2H → 各设备 cudaDeviceSynchronize 或事件同步；主机用锁页内存以支持多设备异步拷贝。

8.3 核心要点回顾

1. 多 GPU 执行：通过 cudaSetDevice 切换当前设备，在每块 GPU 上分配资源并启动核函数；流和事件与创建时所在设备绑定。
2. 点对点：启用 P2P 后可在设备间直接传输或直接访问对方内存；使用 cudaMemcpyPeer 或 UVA 下的 cudaMemcpy 实现高效设备间拷贝。
3. 跨设备同步：用 cudaStreamWaitEvent(streamB, eventA, 0) 让设备 B 的流等待设备 A 上的事件，是实现多 GPU 依赖关系的核心 API。
4. 多设备编程模式：每设备独立内存与流、数据划分、异步传输与计算、必要时 P2P 或主机中转；集群场景使用 CUDA-aware MPI 与 GPUDirect RDMA；结合 CPU-GPU 亲和性可获得更稳定性能。

---

下一章预告

在下一篇博客中，我们将进入第 10 章（全书最后一章），对 CUDA 学习进行总结与拓展：

• 全书知识体系回顾与 CUDA 编程最佳实践
• 性能优化路线图与常见陷阱
• 进阶方向：多流、多 GPU、库与指令、集群与领域库

从多 GPU 与集群，回到整体视角——我们下一章见。

---

本章自测

1. 多 GPU 下「当前设备」的含义是什么？不调用 cudaSetDevice 时默认是哪个设备？
2. 点对点（P2P）传输与经主机中转的 D2H+H2D 相比有何优势？
3. 设备 B 上的流要等待设备 A 上的事件完成，应使用哪个 API？

答案与解析

1. 当前设备指后续的 cudaMalloc、kernel 启动等操作所针对的设备。不调用 cudaSetDevice 时默认为设备 0。
2. P2P 可在支持拓扑下直接在设备内存间传输，不经过主机内存，减少一次 D2H 和一次 H2D，延迟更低、带宽更高；经主机中转需两次 PCIe 传输且占用主机内存。
3. 使用 cudaStreamWaitEvent(stream_B, event_A, 0)，使设备 B 上的 stream_B 等待设备 A 上记录的 event_A 完成后再执行后续操作。

---

系列导航：导读 | 上一篇：第8章 GPU加速库和OpenACC | 下一篇：第10章 程序实现的注意事项

---

本文为「CUDA C编程权威指南」系列博客第 9 篇，共 10 章。基于《Professional CUDA C Programming》by John Cheng, Max Grossman, Ty McKercher。