algorithm - CUDA流压缩算法

Question

我正在尝试使用 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?

Answer 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 。

⁽¹⁾ 严格来说，完全不是传统意义上的流压缩，因为流压缩传统上是一种稳定的算法，但您的要求不包括稳定性。这种宽松的要求可能会导致更有效的实现？