4. 多GPU内存分配与传输

多GPU程序的第一道性能门槛往往不在计算核函数，而在内存。单卡时代我们习惯用 cudaMalloc 申请、cudaMemcpy 搬运，到了多卡环境，同一份数据可能分布在多块 GPU 的显存里，还要在主机与设备、设备与设备之间来回流动。如果对内存类型和传输路径理解不清，很容易写出“算得很快、搬得很慢”的程序。本节围绕显存分配、跨设备拷贝、锁页内存、零拷贝、内存池与异步传输六个要点展开，重点说明它们各自解决什么问题、在什么场景下使用，以及如何组合成一条高效的数据通路。

一、设备内存分配与跨设备拷贝

每块 GPU 拥有独立的显存地址空间，cudaMalloc 在当前设备上分配线性内存，返回的指针只在该设备上下文内有效。多卡程序必须先 cudaSetDevice(i) 再分配，否则会全部落在默认设备 0 上。一个常见的错误是：在设备 0 上分配了缓冲区，却把它当作设备 1 的指针传给核函数，结果要么报非法访问，要么读到错误数据。

跨设备拷贝使用 cudaMemcpyPeer 或带 cudaMemcpyDeviceToDevice 类型的 cudaMemcpy，并显式指定源设备和目标设备。以两块 GPU 为例，把设备 0 上的 d_a 复制到设备 1 的 d_b：

cudaSetDevice(0);
float *d_a; cudaMalloc(&d_a, N * sizeof(float));

cudaSetDevice(1);
float *d_b; cudaMalloc(&d_b, N * sizeof(float));

// 从设备0拷贝到设备1
cudaMemcpyPeer(d_b, 1, d_a, 0, N * sizeof(float));

这条路径是否高效，取决于两块 GPU 之间是否存在 NVLink 或 PCIe P2P 支持。可以用 cudaDeviceCanAccessPeer(&can, 1, 0) 查询，若返回 1，再调用 cudaDeviceEnablePeerAccess(1, 0) 开启对等访问。开启后，设备 1 的核函数可以直接读写设备 0 的显存，省去显式拷贝。若 P2P 不可用，数据只能经主机内存中转，带宽会明显下降。

二、锁页内存与零拷贝内存

普通 malloc 得到的是可分页内存，驱动在传输前需要先把它锁页并复制到临时缓冲区，等于多了一次拷贝。改用 cudaMallocHost 分配锁页内存，DMA 引擎可以直接访问，主机到设备的带宽通常能提升一倍以上。代价是锁页内存不可换出，分配过多会拖慢整个系统，因此一般只对频繁传输的缓冲区使用。

零拷贝内存通过 cudaHostAlloc 配合 cudaHostAllocMapped 标志分配，再用 cudaHostGetDevicePointer 取得设备侧指针。核函数可以直接访问这块主机内存，无需显式拷贝。它适合数据量小、访问次数少、或设备显存不足的场景；但由于每次访问都要走 PCIe，延迟远高于显存，不适合高频读写。

三、内存池与异步传输

频繁 cudaMalloc/cudaFree 会带来明显的同步开销。CUDA 提供流有序内存池：用 cudaMallocAsync 从池中分配、cudaFreeAsync 归还，分配和释放都绑定到具体流，可与计算重叠。相比传统分配，内存池能显著降低反复申请释放的延迟，尤其适合每轮迭代都要临时缓冲区的场景。

异步传输的核心是把拷贝放进流里，用 cudaMemcpyAsync 配合锁页内存，让拷贝与核函数执行重叠。典型模式是“双缓冲”：把数据分成若干块，第 k 块计算的同时，第 k+1 块正在传输。这样 PCIe 传输时间就能被计算掩盖。需要注意的是，异步拷贝的源或目标若涉及主机内存，必须是锁页内存，否则调用会退化为同步。

四、组合使用与判断依据

把上述机制串起来，一条典型的多卡数据通路是：主机用锁页内存准备数据，通过 cudaMemcpyAsync 分块送入各 GPU；GPU 之间若支持 P2P，用 cudaMemcpyPeer 或直接对等访问交换边界数据；临时缓冲区用内存池分配，避免反复申请释放；所有拷贝都挂在流上，与计算重叠。

判断是否用对了，可以观察几个信号：如果 cudaMemcpyAsync 的耗时接近同步拷贝，多半是主机内存没锁页；如果跨设备拷贝带宽远低于 PCIe 理论值，检查 P2P 是否开启；如果 profiler 显示拷贝与计算串行，说明流划分或双缓冲没生效。下面这张检查表可作为快速排查参考：

最后强调一点：多GPU内存管理的目标不是“把数据搬过去”这么简单，而是让数据在正确的时间出现在正确的设备上，并且搬运过程尽量不阻塞计算。理解每条路径的带宽、延迟和适用条件，才能在实际工程中做出合理取舍。