消除内核调用之间的 cudaMemcpy答案

【问题标题】：Eliminate cudaMemcpy between kernel calls消除内核调用之间的 cudaMemcpy
【发布时间】：2015-11-24 00:06:28
【问题描述】：

我有一个被多次调用的 CUDA 内核（100 万不是上限）。我们是否再次启动内核取决于我们的内核返回的标志（result_found）。

for(int i = 0; i < 1000000 /* for example */; ++i) {
    kernel<<<blocks, threads>>>( /*...*/, dev_result_found);
    cudaMemcpy(&result_found, dev_result_found, sizeof(bool), cudaMemcpyDeviceToHost);
    if(result_found) {
        break;
    }
}

分析器说 cudaMemcpy 的执行时间比实际的内核调用要多得多（cudaMemcpy：~88us，cudaLaunch：~17us）。

所以，问题是：

1) 有什么办法可以避免在这里调用 cudaMemcpy？

2) 到底为什么这么慢？将参数传递给内核（cudaSetupArgument）似乎非常快（~0.8 us），而返回结果却很慢。如果我删除 cudaMemcpy，我的程序完成速度会快很多，所以我认为这不是因为同步问题。

【问题讨论】：

还有其他方法（例如零拷贝），但在读取result_found 的值之前，它们仍将依赖于同步（强制内核完成）。您的分析方法可能存在缺陷。 cudaLaunch（运行时 API 调用）的持续时间与内核的持续时间不同。您实际上应该在分析器中查看内核的持续时间（取决于您使用的分析器）。因此，这里的答案很可能只是您误解了探查器数据。但是如果没有更多关于你是如何分析的信息，我不能说得清楚。

标签： cuda

【解决方案1】：

1) 有什么办法可以避免在这里调用 cudaMemcpy？

是的。在这种情况下，动态并行可能会有所帮助。如果您的设备支持它，您可以将i 上的整个循环移到 GPU 上，并从 GPU 启动更多内核。然后启动线程可以直接读取 dev_result_found 并在完成时返回。这完全删除了cudaMemcpy。

另一种方法是大大减少cudaMemcpy 调用的数量。在每个内核启动开始时检查dev_result_found。如果是真的，返回。这样，您只需要在每个 x 迭代中执行一次 memcpy。虽然您会推出比您需要的更多的内核，但这些内核会非常便宜，因为多余的内核会立即返回。

我怀疑这两种方法的组合会提供最佳性能。

2) 到底为什么这么慢？

很难说。我建议您的数字可能有点可疑 - 我猜您正在使用分析器中的 API 跟踪。这会测量 CPU 所看到的时间，因此如果您启动异步调用（内核启动），然后启动同步调用 (cudaMemcpy)，则将使用 memcpy 测量同步成本。

不过，如果您的内核运行速度相对较快，则复制的开销可能会很大。您也无法隐藏任何启动开销，因为您无法异步安排下一次启动。

【讨论】：