将 CUDA cudaMemcpy 分成块

Question

我和一位同事正在集思广益，讨论如何减少主机和设备之间的内存传输时间，结果可能将事情安排到一次超大传输（即一次调用）可能会有所帮助。这导致我创建了一个测试用例，我在其中计算了传输少量大数据块与许多小数据数据块的时间。我得到了一些非常有趣/奇怪的结果，想知道这里有没有人解释一下？

我不会把我的整个代码放在这里，因为它很长，但我用两种不同的方式测试了分块：

1. 显式写出所有 cudaMemcpy，例如：

   cudaEventRecord(开始, 0);
   cudaMemcpy(aD, a, nBytes/10, cudaMemcpyHostToDevice);
   cudaMemcpy(aD + 1*nBytes/10, a + 1*nBytes/10, nBytes/10, cudaMemcpyHostToDevice);
   cudaMemcpy(aD + 2*nBytes/10, a + 2*nBytes/10, nBytes/10, cudaMemcpyHostToDevice);
   cudaMemcpy(aD + 3*nBytes/10, a + 3*nBytes/10, nBytes/10, cudaMemcpyHostToDevice);
   cudaMemcpy(aD + 4*nBytes/10, a + 4*nBytes/10, nBytes/10, cudaMemcpyHostToDevice);
   cudaMemcpy(aD + 5*nBytes/10, a + 5*nBytes/10, nBytes/10, cudaMemcpyHostToDevice);
   cudaMemcpy(aD + 6*nBytes/10, a + 6*nBytes/10, nBytes/10, cudaMemcpyHostToDevice);
   cudaMemcpy(aD + 7*nBytes/10, a + 7*nBytes/10, nBytes/10, cudaMemcpyHostToDevice);
   cudaMemcpy(aD + 8*nBytes/10, a + 8*nBytes/10, nBytes/10, cudaMemcpyHostToDevice);
   cudaMemcpy(aD + 9*nBytes/10, a + 9*nBytes/10, nBytes/10, cudaMemcpyHostToDevice);
   cudaEventRecord(停止, 0);
   cudaEventSynchronize(停止);
   cudaEventElapsedTime(&time, start, stop);

2. 将 cudaMemcpy 放入 for 循环：

   cudaEventRecord(开始, 0);
   for(int i = 0; i {
   cudaMemcpy(aD + i*nBytes/nChunks, a + i*nBytes/nChunks, nBytes/nChunks, cudaMemcpyHostToDevice);
   }
   cudaEventRecord(停止, 0);
   cudaEventSynchronize(停止);
   cudaEventElapsedTime(&time, start, stop);

需要注意的是，为了以防万一，我还在每次测试开始时进行了“热身”传输，尽管我认为不需要它（上下文是由 cudaMalloc 调用创建的）。

我在从 1 MB 到 1 GB 的总传输大小上对此进行了测试，其中每个测试用例传输相同数量的信息，而不管它是如何分块的。我的输出示例如下：

单次大传输 = 0.451616 ms
10 次显式传输 = 0.198016 毫秒
100 次显式传输 = 0.691712 毫秒
10 次循环传输 = 0.174848 毫秒
100 次循环传输 = 0.683744 毫秒
1000 次循环传输 = 6.145792 毫秒
10000 次循环传输 = 104.981247 毫秒
100000 次循环传输 = 13097.441406 毫秒

这里有趣但我没有得到的是，总体而言，10 次传输总是比其他任何传输都快很多，即使是单次大传输！无论数据集有多大或多小，该结果都保持一致（即 10x100MB 与 1x1GB 或 10x1MB 与 1x10MB 仍然会导致 10 倍更快）。如果有人对为什么会这样或者我可能做错了什么得到这些奇怪的数字有任何见解，我很想听听你要说什么。

谢谢！

附：我知道 cudaMemcpy 带有隐式同步，所以我可以使用 CPU 计时器，而 cudaEventSynchronize 是多余的，但我认为最好还是安全一点

更新：我编写了一个函数来尝试利用性能时空连续体中这种明显的裂痕。但是，当我使用该函数时，该函数在我的测试用例中编写为 EXACLTY，效果消失了，我看到了我所期望的（单个 cudaMemcpy 是最快的）。也许这更像是量子物理学而不是相对论，其中观察行为会改变行为......

Answer 1

cudaMemcpy() 是同步的 - CUDA 等到 memcpy 完成后再返回到您的应用程序。

如果您调用 cudaMemcpyAsync()，驱动程序将在 GPU 必须执行 memcpy 之前将控制权返回给您的应用程序。

调用 cudaMemcpyAsync() 而不是 cudaMemcpy() 至关重要。不是因为您希望将传输与 GPU 处理重叠，而是因为这是您获得 CPU/GPU 并发的唯一方法。

在 Amazon EC2 中的 cg1.4xlarge 实例上，驱动程序需要大约 4 微秒来请求 GPU 的内存；因此 CPU/GPU 并发是隐藏驱动程序开销的好方法。

对于您在 10 时看到的差异，我没有现成的解释 - 我希望看到的主要膝盖是 memcpy 大小超过 64K 的位置。驱动程序将小于 64K 的 memcpy 内联到用于提交命令的同一缓冲区中。

Answer 2

在每次cuda调用之前和之后使用cudaThreadSynchronize()来获取真正的内存传输时间，cudaMemcpy()是同步的但不与CPU执行，它取决于调用的函数。

Cuda 函数调用与其他 cuda 函数调用同步，例如其他内存传输或内核执行，这是在 CUDA 开发人员不可见的不同 CUDA 线程中管理的。 cudaMemcpyAsync() 与其他 CUDA 调用异步，这就是为什么它需要复制的 GPU 内存段不与其他并发内存传输重叠。

您确定在这种情况下，在 CUDA 执行线程中同步的 cudaMemcpy() 也与 CPU 线程同步吗？取决于cuda函数可以或不可以，但是如果你在测量时间时使用cudaThreadSynchronize函数，它肯定会与CPU同步，并且每个步骤的实时时间都会出现。

Answer 3

也许这是 CUDA 测量时间的一些特殊性。您正在测量小于 1 毫秒的时间，这非常小。 您是否尝试使用基于 CPU 的计时器对其进行计时并比较结果？