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

第三行是同步的。第四行和第五行是异步的，并且机器支持并发执行。因此，在某个时刻，这两个内核都在GPU上运行。(有可能kernelB在kernelA完成之前开始并完成)。在此期间，主机正在执行第6行。第6行与复制操作同步，但没有阻止它在kernelA或kernelB完成之前执行。

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

- Doug

---

提交到GPU的工作会被提交到流中执行。每个工作流都将按顺序执行。代码示例使用默认流，因此所有操作将按顺序在GPU上执行。CPU可能会在kernelA开始在GPU上执行之前开始执行cudaMemcpy。但是，内存操作将仅在kernelB完成后发生。有关更多信息，请参见CUDA C/C++ Streams and Concurrency网络研讨会。 - Greg Smith

@Doug，你似乎把CPU/GPU并发（异步）和GPU/GPU并发混淆了。在你的例子中，kernelB没有可能在kernelA完成之前启动或完成，因为它们都在NULL流中启动，因此被序列化。如果它们在不同的流中，它们可能会在GPU上并发执行；但是内核调用是异步的（CPU在运行时继续执行）。根据定义，cudaMemcpy()调用是同步的-它在返回之前等待memcpy完成。请参见第2.5节：http://bit.ly/TbZcq4 - ArchaeaSoftware

Answer 1

是的，GPU上有各种队列，驱动程序会管理这些队列。异步调用立即返回，同步调用则需要等待操作完成才能返回。内核调用是异步的，大多数其他CUDA运行时API调用在异步情况下使用后缀Async进行标识。 因此，回答您的问题：

1）GPU中是否有内核队列？（GPU是否会阻塞/停顿主机？）

有各种队列。在同步调用中，GPU会阻塞/停顿主机，但内核启动不是同步操作。 它会立即返回，甚至在内核完成之前或者开始之前就已经返回。将操作启动到单个流中时，该流中的所有CUDA操作都是串行化的。因此，即使内核启动是异步的，如果两个内核在同一个流中启动，你将不会看到它们之间的重叠执行，因为CUDA子系统保证在同一流中的给定CUDA操作不会开始直到先前的所有CUDA操作完成。对于空流（如果您在代码中没有明确调用流，则默认使用的流），还有其他特定规则，但上述描述足以理解这个问题。

2）主机如何知道内核何时完成，并且从设备安全地传输结果到主机？

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

- Robert Crovella

Answer 2

在调用内核后，您可以使用{{link1：cudaDeviceSynchronize（）}}来确保设备完成了所有先前请求的任务。 如果kernelB的结果与kernelA的结果无关，则可以在内存复制操作之前设置此函数。否则，您需要在调用kernelB之前阻止设备，从而导致两个阻塞操作。

- Ian Medeiros

---

从ConcurrencyWebinar中：由于代码示例使用默认流（0），因此在每个CUDA操作之前和之后都有一个隐式的cudaDeviceSynchronize()操作：{cudaMalloc，cudaMemcpy（H->D），cudaMemcpy（D->H）}，即在cudaMemcpy()之前或之后添加cudaDeviceSynchronize()不会产生任何效果。 - Doug