cuda中的内存计算重叠问题答案

【问题标题】：Memory-compute overlap issue in cudacuda中的内存计算重叠问题
【发布时间】：2014-01-28 23:06:19
【问题描述】：

我有一个处理大量数据的 CUDA 内核。由于我无法一次传输所有数据，我必须将它们分成块并逐块处理它们并更新 GPU 上的输出。我正在解析文件中的输入数据。我在想是否可以通过在主机和 GPU 中都有两个缓冲区来重叠块的内存传输。在处理一个卡盘时，我可以读取另一个卡盘，将其传输到 GPU 并将内核启动到同一流。我的问题是内核的执行时间比解析数据并将它们传输到 GPU 慢。鉴于 memcpys 没有阻塞，我如何确保 memcpys 不会覆盖内核使用的数据？

//e.g. Pseudocode
//for every chunk
   //parse data
   //cudaMemcpyAsync ( dev, host, size, H2D )
   //launch kernel
   //switch_buffer
//copy result from device to host

提前谢谢你。

【问题讨论】：

如果你可以选择不要每次都重用相同的内存并切换缓冲区，你可以看看这个webinar。一种可能性是为每对对应的 memcpy 和内核启动使用不同的流。此外，您必须确保使用固定内存！在您的伪代码中，cudaMemcpyAsync 中也缺少使用的 cuda 流。

标签： c++ asynchronous cuda parallel-processing overlap

【解决方案1】：

只需在内核启动后插入带有cudaDeviceSynchronize() 的显式同步点。

这样，您实际上是在同时启动内存传输和启动内核。传输将转到一个缓冲区，内核将在另一个缓冲区上工作。 cudaDeviceSynchronize() 将等到两者都完成，此时您将交换缓冲区并重复。

当然，您还需要在循环内将结果从设备复制到主机，并添加逻辑来处理第一次迭代，当没有数据供内核处理时，以及最后一次迭代，当没有要复制更多数据，但仍要处理一个缓冲区。这可以通过循环内的逻辑或部分展开循环来完成，以专门对第一次和最后一次迭代进行编码。

编辑：

通过将同步点移动到 cudaMemcpyAsync() 之前以及文件读取和解析之后，您允许内核也可以重叠该部分处理（如果内核运行时间足够长的话）。

【讨论】：

我已经用 cudaDeviceSynchronize 实现了这个，但我想避免它，因为它需要时间。尽管如此，我看不到任何其他方法来确保我想要什么。谢谢。
@user2117698，如果你像 Roger 描述的那样双倍缓冲，那么它不会花费太多额外的时间。您的潜在加速受限于您的工作负载在传输和计算之间的平均分布程度。如果两个缓冲区足够大，那么启动内核启动和 memcpy 的开销将被隐藏，性能最终会达到其传输/处理有界的自然程度。