我的应用程序在device-code
中执行一些操作,并在kernel
内生成一个数组。
我需要搜索该数组中第一次出现的元素。我如何在 GPU 中执行它?如果我将数组复制到CPU并在那里进行工作,它将产生大量的内存流量,因为这段代码被调用了很多次。
最佳答案
很可能有一个更复杂的解决方案,但首先,特别是如果元素出现的次数非常少,一个简单的暴力原子最小值可能是一个可行的解决方案:
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;
然后对该索引数组执行“经典”最小归约以获得第一个匹配索引。
关于search - CUDA 中的首次出现搜索,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/17046307/