仅为特定内核调用 cudaDeviceSynchronize()答案

【问题标题】：Calling cudaDeviceSynchronize() only for a particular kernel仅为特定内核调用 cudaDeviceSynchronize()
【发布时间】：2014-08-01 05:21:22
【问题描述】：

我异步调用内核KerA 和KerB。内核KerC 依赖于KerB 整理但独立于KerA。那么如何调用cudaDeviceSynchronize() 以使KerC 等待KerB 完成而不是KerA？

Time -------------------------->
| KerA ------------------------>
| KerB ------> | KerC --------->

【问题讨论】：

标签： asynchronous cuda gpgpu

【解决方案1】：

您可以使用 CUDA 流实现此目的。

如果您不使用任何流，则使用默认流（也称为流 '0'），并且您不会获得并发（就像在每个 CUDA 操作之前和之后插入cudaDeviceSynchronize()，参见these slides） .

但是，如果 KerA 在流 0 中运行，KerB 和 KerC 在流 1 中运行，您将得到您想要的，即 KerB 和 KerC 之间的同步调用，它们是异步的 w.r.t。 KerA。您可以使用cudaStreamSynchronize(streamid) 同步 w.r.t。一个特定的流。

Time ------------------------------------>
| Stream 0: KerA ------------------------>
| Stream 1: KerB ------> | KerC --------->

我链接的slides 中提供了示例。您也可以查看 SDK 的 simpleStreams 或 concurrentKernels 示例。

【讨论】：

谢谢。我以前没有使用过固定内存，所以请你也澄清一下使用“cudaMemcpyAsync with host from 'pinned' memory”的要求。我目前正在使用 Thrust 在我的代码开始时将 100MB 以上的数据上传到没有流的设备 - 如果我稍后在我的代码中使用流，每个流仍然可以读取这些数据吗？我可以使用多少固定内存 - 超过 100MB？如果没有固定内存，流会不会工作？
另外，流可以进一步分成子流吗？
如果你想做一些异步 memcpy，你需要担心固定内存和cudaMallocHost。是的，每个流都可以访问全局内存，但您需要确保不同的流不会产生 RAW/WAW 危害（“并发操作使用的数据应该是独立的”）。不，据我所知，没有子流之类的东西（除非我们谈论的是我想的最新的支持动态并行的 GPU）。