2

我的应用程序在其中做一些事情device-code并在kernel.

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

4

2 回答 2

2

很可能有一个更复杂的解决方案,但首先,特别是如果元素的出现次数非常少,一个简单的暴力 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的元素等于在该索引处搜索的元素,并且如果它不:idxidxINT_MAX

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

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

于 2013-06-11T14:38:16.993 回答
0

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

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

找到第一次出现的另一种方法是使用一种称为并行归约的方法。CUDA SDK 中有一个并行求和的示例(该示例计算数组中所有值的总和)。并行缩减是一个不错的选择,特别是如果您使用具有较旧计算能力版本的硬件并且需要高精度。

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

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

Stack Overflow 上也讨论了原子与减少问题。

于 2013-06-12T17:00:10.693 回答