c++ - OpenCL:循环内核？

Question

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

在我的测试中，调用 clEnqueueNDRangeKernel 的成本很高。例如，当运行模拟的 1000 个子步骤时(需要对 clEnqueueNDRangeKernel 进行 1000 次相同的调用以处理相同的数据)，似乎对 clEnqueueNDRangeKernel 的那些调用实际上成为了瓶颈。我的(伪)代码如下所示:

[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]

据我所知，第一次调用 clEnqeueuNDRangeKernel 会将任何延迟的缓冲区传输初始化到 GPU...因此第一次调用可能会产生额外费用。然而，在我的测试中，10 次迭代的循环比 1000 次迭代快得多，这让我相信数据传输不是瓶颈。

我也觉得 clEnqueueNDRangeKernel 是非阻塞的，因为它在内核完成之前不会阻塞，所以内核的复杂性不应该是瓶颈(在我的例子中，内核在调用 clFinish() 之前不应阻塞执行。

然而，当我分析我的代码时，大部分时间只花在处理 for 循环上，在调用 clFinish() 之前...所以看起来内核本身的排队是花费最多时间的在这里。

我的问题:有没有办法告诉 GPU 重新运行先前排队的内核 N 次，而不是必须手动将内核排队 N 次？在我的情况下，每次迭代都不需要更改或更新内核的参数......内核只需要重新运行。是否可以提高重复调用它的效率？

Answer 1

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

在 OpenCL 1.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 桥(riser 的高延迟？)导致瓶颈，那么 OpenCL2.x 动态并行性必须优于 CPU 端生产者-消费者模式。