为什么在指定的流中没有发生 cudaMemcpyAsync P2P 复制

Question

正如这个问题的答案中所指出的那样

What device number should I use (0 or 1), to copy P2P (GPU0->GPU1)?

源或目标 GPU 上的 cuda 流均可用于点对点复制。然而，这是我从分析中得到的，有点令人困惑。

#include <cuda.h>
#include <cuda_runtime.h>

int main() {
  cudaDeviceEnablePeerAccess(0, 1);
  cudaDeviceEnablePeerAccess(1, 0);

  // on device 0
  cudaSetDevice(0);
  float* data0;
  cudaMalloc(&data0, 1024000);

  // on device 1
  cudaSetDevice(1);
  cudaStream_t stream1;
  cudaStreamCreate(&stream1);
  float* data1;
  cudaMalloc(&data1, 1024000);

  cudaMemcpyAsync(data0, data1, 1024000, cudaMemcpyDefault, stream1);
  cudaDeviceSynchronize();
  cudaMemcpyAsync(data1, data0, 1024000, cudaMemcpyDefault, stream1);
  cudaDeviceSynchronize();
}

问题是，

当我明确地将它分配给设备 1 时，为什么复制作业在设备 0 上执行？

Answer 1

探查器中的索引方案与 cudaSetDevice 函数使用的索引不同。如果您查看您的代码，您创建的流是在设备 1 上创建的，但它与分析器上索引为 [0] 的设备相关联。您应该使用 NVTX's Resource Naming API 来命名设备。它会让您更好地了解资源是如何映射到探查器上的。

Answer 2

这段代码真是粗心！

cudaDeviceEnablePeerAccess 不接受两个设备 ID，第二个参数应始终为 0。

正确调用cudaDeviceEnablePeerAccess后，profiling结果显示无论设置哪个stream，源设备总是会出现cudaMemcpyAsync。似乎如果使用目标 GPU 中的流，它将被重新分配到源设备中的流，这是自动创建的。

不过，虽然复制总是在源设备中执行，但如果将函数下发到目标流，它会阻塞目标流，直到复制完成，我猜是通过cudaEvent。我想这就是为什么

"performance is maximized when stream belongs to the source GPU." as stated in this post What device number should I use (0 or 1), to copy P2P (GPU0->GPU1)?

例如：

cudaMemcpyPeerAsync(data0, 0, data1, 1, 1024000, stream0);
cudaMemcpyPeerAsync(data1, 1, data0, 0, 1024000, stream0);

虽然第一个副本将分配给设备 0，第二个副本分配给设备 1，但这两个操作不会重叠，因为它们在同一个流中。

如果第一个副本设置为设备1，则它们可以重叠。

cudaMemcpyPeerAsync(data0, 0, data1, 1, 1024000, stream1);
cudaMemcpyPeerAsync(data1, 1, data0, 0, 1024000, stream0);

为什么在指定的流中没有发生 cudaMemcpyAsync P2P 复制

Why cudaMemcpyAsync P2P copy is not happening in the specified stream

cuda

gpu

nvidia