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

Question

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

我在找的是一个固定大小的列表，最多有1024个项目的汉明距离（popcount）小于给定的数字。列表顺序并不重要，所以也许我应该要求这些2 ** 20项目的子集。

Answer 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().

Answer 2

OpenCL不支持这样的列表。 OpenCL是一种标准C，有一些扩展和一些限制。你只能操作缓冲区（又名数组）。

您可能要查找的是在运行内核之前需要分配的全局内存缓冲区。在这里，你可以把你的结果放在一个clEnqueueReadBuffer中，你可以检索你的结果。

Answer 3

由于输出预计会比输入要小得多，因此通过原子访问在输出中使用全局索引不会太无效。你需要传递一个包含单个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; 
    } 
}