Cuda：内核启动队列

Question

我没有找到太多关于内核启动操作机制的信息。 API 说看CudaProgGuide。而且我在那里也找不到太多东西。
由于内核执行是异步的，并且某些机器支持并发执行，因此我相信内核有一个队列。

    Host code:      
    1. malloc(hostArry, ......);  
    2. cudaMalloc(deviceArry, .....);  
    3. cudaMemcpy(deviceArry, hostArry, ... hostToDevice);
    4. kernelA<<<1,300>>>(int, int);  
    5. kernelB<<<10,2>>>(float, int));  
    6. cudaMemcpy(hostArry, deviceArry, ... deviceToHost);  
    7. cudaFree(deviceArry);

第 3 行是同步的。第 4 & 5 行是异步的，机器支持并发执行。所以在某些时候，这两个内核都在 GPU 上运行。 （kernelB 有可能在 kernelA 完成之前启动并完成。）此时，主机正在执行第 6 行。第 6 行相对于复制操作是同步的，但没有什么阻止它在 kernelA 或之前执行kernelB 已完成。

1) GPU 中有内核队列吗？ （GPU 会阻塞/停止主机吗？）
2) 主机如何知道内核已经完成，并且将结果从设备 Xfer 到主机是“安全的”？

Answer 1

是的，GPU 上有各种各样的队列，由驱动程序管理。

异步调用或多或少会立即返回。在操作完成之前，同步调用不会返回。内核调用是异步的。如果它们是异步的，大多数其他 CUDA 运行时 API 调用由后缀 Async 指定。所以回答你的问题：

1) GPU 中有内核队列吗？ （GPU 会阻塞/停止主机吗？）

有各种队列。 GPU 在同步调用上阻塞/停止主机，但内核启动不是同步操作。它立即返回，在内核完成之前，甚至可能在内核启动之前。将操作启动到单个流中时，该流中的所有 CUDA 操作都序列化。因此，即使内核启动是异步的，您也不会观察到启动到同一个流的两个内核的重叠执行，因为 CUDA 子系统保证流中的给定 CUDA 操作直到所有先前的 CUDA 操作在相同的流 已经结束。对于空流（如果您没有在代码中显式调用流，您正在使用的流）还有其他特定规则，但前面的描述足以理解这个问题。

2) 主机如何知道内核已经完成，并且将结果从设备 Xfer 到主机是“安全的”？

由于将结果从设备传输到主机的操作是CUDA调用（cudaMemcpy...），并且它与前面的操作在同一流中发出，因此设备和CUDA 驱动程序 管理 cuda 调用的执行顺序，以便 cudaMemcpy 直到所有先前对同一流发出的 CUDA 调用都完成后才开始。因此，在同一流中的内核调用之后发出的cudaMemcpy 保证在内核调用完成之前不会启动，即使您使用cudaMemcpyAsync。

Answer 2

您可以在内核调用之后使用cudaDeviceSynchronize() 来保证之前向设备请求的所有任务都已完成。 如果kernelB 的结果与kernelA 的结果相互独立，则可以在内存复制操作之前设置该函数。如果没有，则需要在调用 kernelB 之前阻塞设备，导致两次阻塞操作。