如何从 OpenCL 内核中获取列表（或子集）？答案

【问题标题】：How to get a list (or subset) from an OpenCL Kernel?如何从 OpenCL 内核中获取列表（或子集）？
【发布时间】：2011-09-20 20:19:07
【问题描述】：

我有一个大数组，上面有 2^20 ulongs。这个小小的 OpenCL 内核就像魅力一样流经它。然而，我完全不知道（谷歌在这里没有帮助）如何从中返回少量项目（2^10）。

我正在寻找的是一个固定大小的列表，其中最多包含 1024 个汉明距离（popcount）小于给定数字的项目。列表顺序无关紧要，所以也许我应该要求这 2**20 个项目的子集。

【问题讨论】：

标签： list return-value return opencl subset

【解决方案1】：

嗯，有一种方法，通过一些技巧。 I forked pyopencl.algorithm and created a new method, sparse_copy_if()，它返回我需要的精确大小的缓冲区，就好像它是一个附加了项目的列表。我会记录下来，然后向 Andreas 提交补丁。

但是，如果您的缓冲区太大，还有一种方法可以进一步提高性能：我按照上面Rick 的建议，创建了一个哈希表，并将所需的结果放入其中。（请注意，总是存在冲突的风险，因此哈希表缓冲区/数组必须比您的预期输出大几个数量级）。

然后，我在哈希表缓冲区上运行 sparse_copy_if() ，只收到一个大小合适的缓冲区。

总结：

我有一个内核扫描 1,000,000 大小的缓冲区。它计算所有结果，但不分离我想要的结果。

然后将这些所需结果放入约 25,000 个缓冲区（哈希表，比原始数据小得多）。

然后，通过在哈希表缓冲区上运行 sparse_copy_if()，您可以获得所需的输出——几乎就像它是一个可以附加项目的列表一样。

当然，

sparse_copy_if() 会产生创建完美大小的缓冲区并将数据复制到其中的开销。但我发现这种开销通常可以弥补，因为您现在正在（低延迟）将小型缓冲区/数组从设备传输回主机。

Code for testing sparse_copy_if() performance versus copy_if().

【讨论】：

【解决方案2】：

由于预计输出远小于输入，因此通过原子访问在输出中使用全局索引不会太无效。您需要传递一个包含单个 uint 的缓冲区，最初设置为 0：

__kernel void K(...,__global uint * outIndex,...)
{
  ...
  if (selected)
  {
    uint index = atomic_inc(outIndex);  // or atom_inc if using OpenCL 1.0 extension
    out[index] = value;
  }
}

【讨论】：

出于兴趣可以写出[atomic_inc(outIndex)] = value;反而？还是编译器不会理解？
优秀的答案！帮助我让我的内核返回结果列表。谢谢。

【解决方案3】：

OpenCL 不支持这样的列表。 OpenCL 是一种具有一些扩展和一些限制的标准 C。您只能对缓冲区（也称为数组）进行操作。

您可能会寻找在运行内核之前需要分配的全局内存缓冲区。在这里，您可以将结果放入并使用 clEnqueueReadBuffer 检索结果。

【讨论】：