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

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

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 桥（提升器的高延迟？）造成瓶颈，那么 OpenCL2.x 动态并行性必须优于 CPU 端生产者-消费者模式。