可以同时完成多少个 cudaMemcpyAsync 操作？

Question

考虑以下情况：

//thread 0 on device 0:
cudaMemcpyAsync(Dst0, Src0, ..., stream0);//stream0 is on Device 0;

...
//thread 1 on device 1:
cudaMemcpyAsync(Dst1, Src1, ..., stream1);//stream1 is on Device 1;

两个 memcpy 操作能否同时发生并获得双倍的主机设备带宽（只要主机内存带宽足够）？如果答案是肯定的，那么这种并发性是否有上限？

我计划为单个计算节点中的许多 (6-8) 个 GPU 编写一些程序，因此这对性能非常关键。

Answer 1

实际上只有一个cudaMemcpy 操作可以在任何给定时间、每个方向上通过 PCIE 链路传输数据。其他操作当然可以排队，但只有一个可以实际使用链接。

一些 CUDA GPU 有一个 DMA 引擎，其他有 2 个。有两个的那些可以同时双向传输数据。这就是数据传输中并发的限制：每个方向一个（假设双复制引擎）。

假设数据传输不小，那么在每次传输期间，PCIE 总线将被充分利用（在那个方向上）。