CUDA 中的首次出现搜索答案

【问题标题】：First occurrence search in CUDACUDA 中的首次出现搜索
【发布时间】：2013-06-11 14:13:54
【问题描述】：

我的应用程序在device-code 中做一些事情并在kernel 中生成一个数组。

我需要搜索该数组中第一次出现的元素。我如何在 GPU 中执行它？如果我将数组复制到 CPU 并在那里进行工作，它会产生如此多的内存流量，因为这段代码被调用了很多次。

【问题讨论】：

【解决方案1】：

很可能有一个更复杂的解决方案，但首先，特别是如果元素的出现次数非常少，一个简单的暴力 atomic-min 可能是一个可行的解决方案：

template<typename T> __global__ void find(T *data, T value, int *min_idx)
{
    int idx = threadIdx.x + blockDim.x*blockIdx.x;
    if(data[idx] == value)
        atomicMin(min_idx, idx);
}

如果出现的次数真的很小，因此几乎所有线程甚至都不尝试访问原子，这实际上可能不是那么糟糕的解决方案。否则（如果搜索到的元素不是那么稀有），您将有更多的经线内发散，更糟糕的是，冲突原子操作的可能性要高得多。

编辑：对于更复杂的方法（但可能仍然不是最好的），您还可以在前一步创建一个 int 数组，并将索引 idx 处的值设置为idx 如果输入数组的元素等于在该索引处搜索的元素，INT_MAX 如果不等于：

indices[idx] = (data[idx]==value) ? idx : INT_MAX;

然后对该索引数组进行“经典”最小缩减以获得第一个匹配索引。

【讨论】：

我会等着看其他人是否有新的想法，否则我会选择你的答案作为最佳答案。谢谢。
请注意：如果您的数组比您拥有的线程数大得多，您可以执行逐步减少（即每个线程只加载一个值，在线程之间进行减少且仅当未找到该值时 - 迭代）

【解决方案2】：

一种方法是使用atomic 操作来阻止其他线程访问可编辑数据，直到当前处理完成。

以下是查找单词第一次出现的示例： http://supercomputingblog.com/cuda/search-algorithm-with-cuda/ 该示例中使用了atomicMin 函数。另外，文中还有GPU和CPU的性能对比。

另一种查找第一次出现的方法是使用一种称为并行归约的方法。 CUDA SDK 中有一个并行求和的示例（该示例计算数组中所有值的总和）。并行缩减是一个不错的选择，尤其是在您使用具有较旧计算能力版本的硬件并且需要高精度时。

要使用并行归约来查找第一次出现，您首先检查数组中的值是否等于您要查找的值。如果是，则保存其索引。然后，您执行一个或多个min 操作（不是原子最小值），在其中比较您在上一步中保存的索引。您可以通过编辑 CUDA SDK 的并行求和示例来实现此搜索。

This site 有一些关于归约和原子操作的信息。它还包括二叉树缩减和解决方法原子函数，我在这里没有谈到。

还讨论了原子与还原问题on Stack Overflow。

【讨论】：

抱歉，可以更好地解释如何将并行归约转化为这种情况，您告诉我这些min 和max 函数是什么？请问，你能从文档中链接它们吗？谢谢你:)