为什么指定流中没有发生 cudaMemcpyAsync P2P 复制答案

【问题标题】：Why cudaMemcpyAsync P2P copy is not happening in the specified stream为什么指定流中没有发生 cudaMemcpyAsync P2P 复制
【发布时间】：2015-03-24 07:03:56
【问题描述】：

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

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();
}

问题是，

当我将设备 0 显式分配给设备 1 时，为什么要在设备 0 上执行复制作业？

【问题讨论】：

我当然可以提供帮助，但您似乎在这里问了 5 个不同的问题。标题中的问题很容易回答，但您似乎正在寻找有关该主题的论文。在一个 SO 问题中提出多个问题并不是一个好主意，它会使问题变得更加难以回答。如果您需要帮助，我建议您提出一个重点问题。你读过this question/answer 吗？
@RobertCrovella 感谢您的指出，我已经检查了您建议的链接，并完善了我的问题。
我同意之前的反馈。请考虑更新您的“问题”文本，以便将其表述为问题。

标签： cuda gpu nvidia

【解决方案1】：

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

【讨论】：

鉴于我在这里提出的问题，这似乎是答案。但是，如果您看到问题stackoverflow.com/posts/29226727/revisions 的先前版本，您会发现情况并非如此。如图所示，无论我在哪个设备上进行复制操作，profile结果总是显示在设备0上。

【解决方案2】：

这段代码真的很粗心！

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

cudaDeviceEnablePeerAccess被正确调用后，分析结果显示cudaMemcpyAsync总是发生在源设备中，无论设置哪个流。似乎如果使用目标 GPU 中的流，它将被重新分配给源设备中的流，该流是自动创建的。

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

“当流属于源 GPU 时性能最大化。”如本文所述 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);

【讨论】：