我正在尝试使用 CUDA 构建一个并行算法,该算法接受一个整数数组并删除所有 0
,保留或不保留顺序。
例子:
全局内存:{0, 0, 0, 0, 14, 0, 0, 17, 0, 0, 0, 0, 13}
主机内存结果:{17, 13, 14, 0, 0, ...}
最简单的方法是使用主机在O(n)
时间内删除0
。但考虑到我有大约 1000
元素,将所有内容留在 GPU 上并在发送之前先压缩它可能会更快。
首选方法是创建一个设备上的堆栈,这样每个线程都可以弹出和推送(以任何顺序)到堆栈上或从堆栈中取出。但是,我不认为 CUDA 有这个的实现。
等效(但慢得多)的方法是继续尝试写入,直到所有线程都完成写入:
kernalRemoveSpacing(int * array, int * outArray, int arraySize) {
if (array[threadId.x] == 0)
return;
for (int i = 0; i < arraySize; i++) {
array = arr[threadId.x];
__threadfence();
// If we were the lucky thread we won!
// kill the thread and continue re-reincarnated in a different thread
if (array[i] == arr[threadId.x])
return;
}
}
此方法的唯一好处在于我们将在 O(f(x))
时间内执行,其中 f(x)
是非零的平均数数组中的值 (f(x) ~= ln(n)
对于我的实现,因此 O(ln(n))
时间,但具有很高的 O
常量)
最后,诸如快速排序或归并排序之类的排序算法也可以解决该问题,并且实际上可以在 O(ln(n))
相对时间内运行。我认为甚至可能有比这更快的算法,因为我们不需要浪费时间排序(交换)零-零元素对和非零非零元素对(不需要保持顺序)。
So I'm not quite sure which method would be the fastest, and I still think there's a better way of handling this. Any suggestions?
最佳答案
您要的是一种称为流压缩1 的经典并行算法。
如果推力是一个选项,您可以简单地使用 thrust::copy_if
.这是一个稳定的算法,它保留了所有元素的相对顺序。
草图:
#include <thrust/copy.h>
template<typename T>
struct is_non_zero {
__host__ __device__
auto operator()(T x) const -> bool {
return x != 0;
}
};
// ... your input and output vectors here
thrust::copy_if(input.begin(), input.end(), output.begin(), is_non_zero<int>());
如果 Thrust 不是一个选项,您可以自己实现流压缩(有很多关于该主题的文献)。这是一个有趣且相当简单的练习,同时也是更复杂的并行原语的基本构建 block 。
(1) 严格来说,完全不是传统意义上的流压缩,因为流压缩传统上是一种稳定的算法,但您的要求不包括稳定性。这种宽松的要求可能会导致更有效的实现?
关于algorithm - CUDA流压缩算法,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/34059753/