OpenCL:循环内核

OpenCL: loop kernel?

本文关键字：内核 循环 OpenCL      更新时间：2023-10-16

我正在运行一个OpenCL内核，它反复处理和重新处理相同的数据集(它是一个迭代物理求解器)。

在我的测试中，调用clEnqueueNDRangeKernel需要付出不小的代价。例如，当运行模拟的1000个子步骤(需要对clEnqueueNDRangeKernel进行1000次相同的调用来处理相同的数据)时，对clEnQueue NDRangekernel的这些调用似乎实际上成为了瓶颈。我的(伪)代码如下：

[create buffers]
[set kernel arguments]
for (int i = 0; i < 1000; i++) //queuing the kernels takes a while
{
clEnqueueNDRangeKernel(queue, kernel, args...); 
}
clFinish(queue); //waiting for the queue to complete doesn't take much time
[read buffers]

我知道对clEnqueuNDRangeKernel的第一次调用将初始化任何延迟的缓冲区传输到GPU。。。因此第一次呼叫可能具有额外的成本。然而，在我的测试中，10次迭代的循环比1000次迭代快得多，这让我相信数据传输不是瓶颈。

我还认为clEnqueueNDRangeKernel是非阻塞的，因为它在内核完成之前不会阻塞，所以内核的复杂性不应该成为瓶颈(在我的情况下，内核执行不应该阻塞，直到调用clFinish())。

然而，当我分析我的代码时，在调用clFinish()之前，大部分时间只用于处理for循环。。。因此，内核本身的排队似乎是这里花费最多时间的。

我的问题：有没有一种方法可以告诉GPU重新运行之前排队的内核N次，而不是手动将内核排队N次？在我的情况下，每次迭代都不需要更改或更新内核的参数。。。内核只需要重新运行即可。反复调用它能提高效率吗？

OpenCL 2.x支持动态并行，允许1个工作项启动新内核。如果每次内核启动都不需要任何gpu cpu数据传输，那么可以让1个工作项启动1000个内核，并等待每个内核在该工作项之前完成。使用事件使所有子内核相继运行。

在OpenCL1.2中，您可以使用原子和循环来进行"运行中线程"内核同步，但这不会比新内核启动imo更快，而且这不是一种可移植的同步方式。

如果每次内核启动的时间比每次内核运行的时间长，那么GPU上的工作就不够了。你可以简单地在GPU上执行c=a+b，这还不够快，因为GPU管道上的内核调度需要比执行c=a+b更多的时间。

但是，您仍然可以使用clEnqueueWaitForEvents和按顺序的命令队列执行以下方法：

thread-0:
enqueue user event, not triggered
enqueue 1000 kernels, they don't start yet because of untriggered wait
thread-1:
nothing

下一个时间步骤：

thread-0:
enqueue new user event on a new command queue, not triggered
enqueue 1000 kernels on new command queue so they don't start yet
thread-1:
run the old command queue from last timestep by triggering the user event

这样排队和跑步至少可以"重叠"。如果你需要更多的排队来运行重叠率，

thread-0 to thread-N-2:
enqueue new user event on a new command queue, not triggered
enqueue 1000 kernels on new command queue so they don't start yet
thread-N-1:
iterate all command queues
run currently selected command queue from last timestep  by triggering the user event

现在您的入队速度提高了N-1倍，运行它们将只需要GPU侧的调度开销。如果你有很多GPU侧调度开销(1M的工作项对于1M的c=a+b计算)，那么你应该为每个工作项做更多的工作。

也许让生产者-消费者风格的内核启动更好，其中7个线程生成填充的命令队列，等待在自己的用户事件上触发，8个线程。线程通过触发它们来消耗它们。即使他们需要将下载数据上传到GPU或从GPU上传数据，这也会起作用。

即使是像HD7870这样的旧GPU也同时支持32+个命令队列(每个GPU)，因此您可以通过高端CPU来提高排队性能。

如果pci-e桥(risers的高延迟？)是造成瓶颈的原因，那么OpenCL2.x动态并行性必须比CPU端生产者-消费者模式更好。

相关文章：