消除内核调用之间的cudaMemcpy

我有一个多次调用的CUDA内核（100万不是限制）。我们是否再次启动内核取决于flag（result_found），我们的内核返回。消除内核调用之间的cudaMemcpy

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，我的程序会更快完成，所以我认为这不是因为同步问题。

来源

2015-11-23 Kirill Kuvshinov

还有其他方法（例如零拷贝），但是在读取'result_found'的值之前，它们仍然依赖于同步（迫使内核完成）。你的分析方法可能有缺陷。 'cudaLaunch'（运行时API调用）的持续时间与内核的持续时间不同。实际上，您应该查看Profiler中内核的持续时间（取决于您使用的是哪个分析器）。所以，这里的答案很可能就是你误解了剖析器数据。但是如果没有更多关于你如何分析的信息，我不能一概而论。 –

1）有没有办法避免在这里调用cudaMemcpy？

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

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

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

2）为什么它毕竟如此缓慢？

很难说。我建议你的号码可能有点可疑 - 我想你正在使用探查器的API跟踪。这将测量CPU的时间，因此如果启动异步调用（内核启动），然后执行同步调用（cudaMemcpy），则会使用memcpy测量同步成本。

尽管如此，如果你的内核运行速度相对较快，副本的开销可能会很大。您也无法隐藏任何启动开销，因为您无法异步安排下一次启动。

来源

2015-11-23 16:06:28 Jez

消除内核调用之间的cudaMemcpy

回答

相关问题