使用cudaMemcpyPeer/Async时需启用单向还是双向CUDA对等访问？

核心问题 #

使用cudaMemcpyPeer或cudaMemcpyPeerAsync实现设备间不经过主机的直接数据拷贝时，需要启用单向还是双向CUDA对等访问？是否需要同时执行两处cudaDeviceEnablePeerAccess调用？另外，调用拷贝接口时的活跃设备是否会影响对等访问的要求？

对等访问启用规则 #

cudaDeviceEnablePeerAccess的作用是让当前设置的活跃设备获得对目标设备内存的访问权限，权限是单向且基于发起访问的设备的。针对不同的拷贝发起场景，对等访问的要求如下：

场景1：活跃设备为源设备（如示例中设备0） #

此时需要在源设备上启用对目标设备的对等访问：

cudaSetDevice(0);
cudaDeviceEnablePeerAccess(1, 0);

因为源设备要将数据写入目标设备的内存空间，必须拥有访问目标设备的权限。

场景2：活跃设备为目标设备（如示例中设备1） #

此时需要在目标设备上启用对源设备的对等访问：

cudaSetDevice(1);
cudaDeviceEnablePeerAccess(0, 0);

因为目标设备要从源设备的内存空间读取数据，必须拥有访问源设备的权限。

场景3：活跃设备为第三方设备（如示例中设备2） #

此时需要在第三方设备上同时启用对源设备和目标设备的对等访问：

cudaSetDevice(2);
cudaDeviceEnablePeerAccess(0, 0);
cudaDeviceEnablePeerAccess(1, 0);

因为第三方设备作为拷贝操作的调度方，需要同时访问源、目标设备的内存才能完成拷贝。

双向访问的必要性 #

不需要强制启用双向对等访问，除非你需要在两个方向都发起拷贝（比如既从设备0拷到1，又从1拷到0）。仅针对单一方向的拷贝，满足上述对应场景的单向访问要求即可。

示例代码对应说明 #

针对示例中的拷贝操作cudaMemcpyPeerAsync(dst1, 1, src0, 0, nbytes);：

• 若活跃设备是0：仅需启用设备0对1的对等访问，原代码中第一处调用足够；
• 若活跃设备是1：仅需启用设备1对0的对等访问，原代码中第二处调用足够；
• 若活跃设备是2：原代码中的两处调用完全无效，必须在设备2上分别启用对0和1的对等访问。

内容的提问来源于stack exchange，提问作者user2052436