2. CUDA多GPU编程模型

单GPU编程时，我们习惯把“主机线程—设备线程—显存”看成一条直线：CPU 发起 kernel，GPU 用 grid/block/thread 组织并行，数据放在这块卡的显存里。一旦系统里出现多块 GPU，这条直线就变成了一个二维甚至三维的结构：主机侧可能有多个线程或进程，设备侧有多个独立的执行上下文，每块 GPU 有自己的显存空间，而它们之间又可能通过 PCIe 或 NVLink 互连。理解多 GPU 编程模型，本质上就是理解“谁在什么上下文里、访问哪块内存、走哪条链路”。

主机线程模型与设备执行上下文

CUDA 的执行上下文（context）是驱动层管理 GPU 资源的基本单位，它绑定一块具体的设备，持有该设备上的模块、内存分配、流和事件等状态。单进程单卡时，运行时 API 会在第一次调用 CUDA 函数时自动创建主上下文（primary context）。多 GPU 场景下，每块被使用的设备都会有一个对应的上下文，主机线程通过 cudaSetDevice(i) 切换“当前设备”，后续的 cudaMalloc、kernel 启动、流操作都作用在这个当前设备上。

这里有一个容易踩坑的点：当前设备是线程局部的。如果主线程设置了 device 0，再创建一个新线程去启动 kernel，新线程的当前设备默认是 0，而不是继承主线程的设置。因此多线程驱动多卡时，每个线程进入工作前都应显式调用 cudaSetDevice。常见做法是“一进程一卡”或“一进程多卡多线程”，前者隔离性最好，后者通信开销更低但需要自己管理线程与设备的绑定关系。

设备侧的线程模型没有因为多卡而改变：kernel 仍然以 grid 为单位启动，grid 内是 block，block 内是 thread，线程通过 blockIdx、threadIdx 定位。变化在于，一个 grid 只能在一台设备上执行，不能跨卡拆分。多卡并行是靠“多个 grid 分别提交到不同设备”实现的，而不是靠一个 grid 内部跨设备调度。

多GPU内存空间与对等访问

每块 GPU 的显存是独立的物理地址空间。默认情况下，device 0 上的指针不能直接在 device 1 的 kernel 里解引用，否则会触发非法访问。要让一块卡访问另一块卡的内存，需要设备间对等访问（Peer-to-Peer，P2P）。P2P 的前提是拓扑支持：同一 PCIe 交换机下的卡、或通过 NVLink 直连的卡通常可以 P2P，跨 NUMA 节点、经过 CPU 根复合体的路径则可能不支持。

启用 P2P 的典型流程是：先用 cudaGetDeviceCount 获取设备数，用 cudaDeviceCanAccessPeer(&can, i, j) 查询 i 能否访问 j，再在 i 的上下文里调用 cudaDeviceEnablePeerAccess(j, 0) 建立单向访问能力。注意它是单向的，i 访问 j 和 j 访问 i 要分别启用。启用后，device i 的 kernel 可以直接读写 device j 上 cudaMalloc 出来的指针，底层由驱动把访问路由到对端显存。

P2P 的收益取决于链路。NVLink 带宽远高于 PCIe，适合频繁的小消息或大块传输；PCIe P2P 虽然也能用，但带宽受限于 PCIe 通道。如果查询发现不支持 P2P，退路是让数据经主机内存中转：先从源卡 cudaMemcpy 到主机，再从主机拷到目标卡，代价是两次拷贝和同步。

统一虚拟地址（UVA）

统一虚拟地址把主机内存和所有设备的显存映射到同一个虚拟地址空间。开启 UVA 后，一个指针的数值本身就携带了“它属于哪块设备”的信息，驱动可以据此判断访问是否需要跨设备路由。这带来两个直接好处：一是可以用同一个指针变量在主机代码里指向不同设备的内存，简化管理；二是 cudaMemcpyDefault 能根据指针自动推断拷贝方向和设备，不必再手写 cudaMemcpyDeviceToDevice 之类的枚举。

需要区分 UVA 和统一内存（Unified Memory）。UVA 只是地址空间统一，内存仍然物理地待在原设备上，访问远端内存依然要走 P2P 或主机中转；统一内存则进一步允许按需迁移页面，让 kernel 像访问本地内存一样访问远端数据，由驱动在缺页时搬移。本节讨论的多卡编程主要建立在 UVA + P2P 之上，统一内存是另一条可选路径。

多GPU执行上下文管理实践

下面这段示例演示了双卡环境下查询 P2P 能力、启用对等访问，并把数据从 device 0 直接拷到 device 1 的过程。参数为示例设定：假设机器有 2 块 GPU，且 cudaDeviceCanAccessPeer 返回 1。

int n = 1 << 20;                 // 每个 float 4 字节，共 4 MB
size_t bytes = n * sizeof(float);

int devCount = 0;
cudaGetDeviceCount(&devCount);    // 示例中 devCount = 2

// 在 device 0 上分配并初始化
cudaSetDevice(0);
float *d0 = nullptr;
cudaMalloc(&d0, bytes);
cudaMemset(d0, 0, bytes);

// 在 device 1 上分配
cudaSetDevice(1);
float *d1 = nullptr;
cudaMalloc(&d1, bytes);

// 查询 0 能否访问 1，并启用
int can = 0;
cudaDeviceCanAccessPeer(&can, 0, 1);
if (can) {
    cudaSetDevice(0);
    cudaDeviceEnablePeerAccess(1, 0);
}

// 从 device 0 直接拷贝到 device 1（UVA 下可用 Default）
cudaSetDevice(0);
cudaMemcpy(d1, d0, bytes, cudaMemcpyDefault);

cudaDeviceSynchronize();

执行后应检查每一步的返回码。若 cudaDeviceCanAccessPeer 返回 0，说明拓扑不支持，此时 cudaDeviceEnablePeerAccess 会返回 cudaErrorPeerAccessUnsupported，程序应回退到主机中转路径。若拷贝成功，cudaMemcpy 返回 cudaSuccess，且 device 1 上的数据与 device 0 一致，可用一次回拷到主机做校验。

上下文、流与同步的关系

多卡并行时，每块设备上的工作由各自的流（stream）组织。默认流是同步语义，跨设备的默认流之间没有隐式同步关系，因此“device 0 的 kernel 完成”并不自动意味着“device 1 的拷贝可以开始”。如果两块卡之间存在数据依赖，必须显式同步：用事件（event）记录源设备上的完成点，再让目标设备等待该事件，或者退一步用 cudaDeviceSynchronize 做粗粒度同步。事件跨设备使用时，要确保记录和等待发生在正确的当前设备上下文中。

上下文管理的另一个实践要点是资源释放顺序。P2P 访问能力是挂在上下文上的，关闭设备或销毁上下文前，应先 cudaDeviceDisablePeerAccess，再 cudaFree 各设备上的内存，最后让进程退出。顺序颠倒可能导致释放失败或残留状态。对于长时间运行的服务，建议把“设备选择—上下文初始化—P2P 启用—内存池建立”封装成一次性初始化流程，运行期只做流提交和事件同步，避免频繁切换当前设备带来的额外开销。

把上面几个概念串起来看：主机线程通过 cudaSetDevice 选择设备，每块设备有独立的执行上下文和显存空间；UVA 让指针自带设备信息，P2P 让跨卡访问不必经过主机；而流和事件负责在多个上下文之间建立正确的先后关系。多 GPU 程序的正确性，往往不取决于 kernel 写得多快，而取决于这些上下文、内存和同步关系有没有理清楚。