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);

第三行是同步的。第4&5行是异步的，机器支持并发执行。所以在某种程度上，这两个内核都在GPU上运行。(在kernelA结束之前，kernelB可能会启动和结束。)当发生这种情况时，主机正在执行第6行。第6行对于复制操作是同步的，但在kernelA或kernelB完成之前，没有任何东西阻止它执行。

1) GPU中是否有内核队列？( GPU是否阻止/停止主机？)

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

Answer 1

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

异步调用或多或少会立即返回。直到操作完成，同步调用才会返回。内核调用是异步的。大多数其他CUDA运行时应用编程接口调用如果是异步的，则由后缀Async指定。所以我来回答你的问题：

1) GPU中是否有内核队列？( GPU是否阻止/停止主机？)

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

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

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

Answer 2

您可以在内核调用之后使用cudaDeviceSynchronize()来保证之前请求到设备的所有任务都已完成。如果kernelB的结果独立于kernelA上的结果，则可以在内存复制操作之前设置此函数。如果不是，则需要在调用kernelB之前阻塞设备，从而导致两个阻塞操作。