如何避免在OpenCL中回读

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

我使用这个数字来决定下一个循环需要多少个工作项目;和
我用这个数字来决定何时退出循环（当数字是0）。

我发现阅读需要大部分时间的循环。有什么办法可以避免它？

一般来说，如果您需要反复调用一个内核，并且退出条件依赖于内核生成的结果（不是固定数字循环），那么您如何有效地执行该操作？有什么像OpenGL中的遮挡查询，你可以做一些查询，而不是从GPU回读？

来源

2012-07-20 redpearl

你有几个选择这里：

如果可能的话，你可以简单地移动循环和条件到内核？使用一种方案，其中额外的工作项目根据当前迭代的输入而不执行任何操作。
如果1.是不可能的，我建议您将由“决定”内核生成的数据存储在缓冲区中，并使用该缓冲区“引导”其他内核。

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

来源

2012-07-20 19:42:59 Ani

如果我启动很多工作项目（比如1000），并且其中大部分（比如说990）都不做任何事情。例如。他们测试一个条件并立即返回，如果它是假的。剂量它有类似的速度，只是启动10个工作项目？还是会慢得多？ – redpearl 2012-07-20 19:46:10

我不明白为什么这会是一件坏事。只要没有回读，它应该只要工作组中最长的工作项目一样长，是正确的？ – Ani 2012-07-20 19:47:31

我的真实程序比较复杂。我在每个循环中调用许多内核。我不认为有可能将整个循环的工作实现到内核中。 :( – redpearl 2012-07-20 19:54:43

从GPU读取数字内核总是需要10s - 1000s微秒或更多。

如果控制编号总是减少，您可以保留在全局内存中，并根据全局标识进行测试，并确定每次迭代时内核是否工作。使用全局内存屏障来同步所有线程...

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-20 20:16:12

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

我们发现了两件事情：

使用两个（或更多）的缓冲区！在内核的第一次迭代中，对b1中的数据进行操作，然后对b2中的下一个操作，然后再对b1中的数据进行操作。在每个内核调用之间的中，回读其他缓冲区的结果并检查是否该停止迭代。内核花费的时间比读取时间长时效果最好。使用一个分析工具来确保你没有等待读取（如果你是，增加缓冲区的数量）。
过度拍摄！为每个内核添加一个完成检查，并在将数据复制回来之前将其称为几（100s）次。如果你的内核是低成本的，这可以很好地工作。

来源

2012-07-23 02:23:08

如何避免在OpenCL中回读

回答

相关问题