3

我正在用 OpenCL 实现一个算法。我将在 C++ 中循环多次，并且每次都调用同一个 OpenCL 内核。内核将生成下一次迭代的输入数据和这些数据的数量。目前，我在每个循环中读回这个数字有两种用途：

  1. 我使用这个数字来决定下一个循环需要多少工作项；和
  2. 我使用这个数字来决定何时退出循环（当数字为 0 时）。

我发现阅读花费了循环的大部分时间。有什么办法可以避免吗？

一般来说，如果你需要重复调​​用一个内核，并且退出条件依赖于内核生成的结果（不是固定数量的循环），你如何高效地做到这一点？有没有像 OpenGL 中的遮挡查询这样的东西，你可以只做一些查询而不是从 GPU 读回？

4

3 回答 3

2

从 GPU 内核读回一个数字总是需要 10s - 1000s 微秒或更长时间。

如果控制数一直在减少，您可以保留在全局内存中，并根据全局 id 对其进行测试，并确定内核在每次迭代时是否正常工作。使用全局内存屏障来同步所有线程......

kernel void x(global int * the_number, constant int max_iterations, ... )
{
    int index = get_global_id(0);
    int count = 0;    // stops an infinite loop

     while( index  < the_number[0] && count < max_iterations )
     {
      count++;
      // loop code follows

      ....

      // Use one thread decide what to do next 
      if ( index  == 0 )
      {
          the_number[0] = ... next value
      }

      barrier( CLK_GLOBAL_MEM_FENCE ); //  Barrier to sync threads
    }
}
于 2012-07-20T20:16:12.257 回答
1

你有几个选择：

  1. 如果可能，您可以简单地将循环和条件移动到内核中吗？根据当前迭代的输入，使用附加工作项不执行任何操作的方案。
  2. 如果 1. 不可能，我建议您将“决策”内核生成的数据存储在缓冲区中，并使用该缓冲区“指导”其他内核。

这两个选项都允许您跳过回读。

于 2012-07-20T19:42:59.310 回答
0

我刚刚完成了一些我们必须解决这个确切问题的研究！

我们发现了几件事：

  1. 使用两个（或更多）缓冲区！让内核的第一次迭代对 b1 中的数据进行操作，然后对 b2 进行下一次迭代，然后再对 b1 进行操作。在每次内核调用之间，读回另一个缓冲区的结果并检查是否该停止迭代。当内核花费的时间比读取时间长时效果最好。使用分析工具确保您没有等待读取（如果是，请增加缓冲区的数量）。

  2. 超拍！为每个内核添加一个完成检查，并在将数据复制回来之前调用它几次（100 秒）。如果你的内核是低成本的，这可以很好地工作。

于 2012-07-23T02:23:08.060 回答