缺省流 , 内核异步的一个小困惑

我看一个 页面的东西

cudaMemcpy(d_a, a, numBytes, cudaMemcpyHostToDevice);
increment<<<1,N>>>(d_a)

myCpuFunction(b)

cudaMemcpy(a, d_a, numBytes, cudaMemcpyDeviceToHost);

他前面说 这些 cudaMemcpy increment cudaMemcpy 都是由 主机发射, 在内核执行的。
当host 发射了increment ,increment 会 立刻返回 (执行在主机), host 继续下边的 myCpuFunction ,从而实现了
increment 和 myCpuFunction 并行。

后边有句说:
Whether the host function or device kernel completes first doesn’t affect the subsequent device-to-host transfer, which will begin only after the kernel completes.
这句说明:内核执行完,就会执行下边的cudaMemcpy ,不管主机是否执行完 myCpuFunction.
但是,我想,主机没有执行完 myCpuFunction , 它就没有发射 cudaMemcpy(a, d_a,…) 怎么可能不管主机是否执行完 myCpuFunction ?

求解释。。。。

LZ您好:

kernel是异步返回的,所以主机端在kernel还没有执行完之前就返回了,这样可以执行后面的myCpuFunction(b)。此时实现了GPU端执行kernel和CPU端执行myCpuFunction(b)的并行。

Whether the host function or device kernel completes first doesn’t affect the subsequent device-to-host transfer, which will begin only after the kernel completes.
这句话,我觉得作者的本意是指,无论kernel先结束还是myCpuFunction(b),都不会造成后面回拷数据在逻辑上的正确性。回拷总是在kernel结束之后才执行。

稍微分析下:
1:如果myCpuFunction(b)的执行时间小于等于kernel的执行时间,我们取一个极限情况,myCpuFunction(b)的执行时间为0 ,这相当于直接kernel后面写数据回拷。因为他们在同一个stream内,后面的回拷总是保证前面kernel执行完毕以后才运行的,这个是自动保证的,无须担心。

2:如果myCpuFunction(b)执行的时间大于kernel的执行时间,那么myCpuFunction(b)执行完毕之后才发布回拷的命令,此时也绝对是在kernel结束以后了。(GPU端可能还空闲等待了若干时间)

所以,回拷总保证是在kernel执行完毕之后才进行。
我觉得作者是这个意思,而不是说回拷在kernel完成之后总是能立即执行。

大致推断如此,仅供参考,您可以向作者进一步求证。

祝您编码顺利~

谢谢斑竹,原来是这样 的。 。 我就说了。。。