CUDA流压缩算法

一尘不染

CUDA流压缩算法

algorithm

我正在尝试使用CUDA构建一个并行算法，该算法采用整数数组并删除所有0带有或不保留顺序的。

例：

全局内存：{0，0，0，0，14，0，0，17，0，0，0，0，13}

主机内存结果：{17、13、14、0、0，…}

最简单的方法是使用主机删除0中的O(n)时间。但是考虑到我周围有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常数）

最后，排序算法（例如quicksort或mergesort）也可以解决该问题，并且实际上是在O(ln(n))相对时间内运行的。我认为甚至有一种算法可以比这更快，因为我们不需要浪费时间排序（交换）零零元素对和非零非零元素对（不需要保持顺序）。

因此，我不太确定哪种方法最快，而且我仍然认为有更好的方法来处理此问题。有什么建议？

阅读 545

2020-07-28

共1个答案

一尘不染

您需要的是一种经典的并行算法，称为 流压缩 1。

如果选择“推力”，则可以简单地使用thrust::copy_if。这是一种稳定的算法，它保留所有元素的相对顺序。

草图：

#include <thrust/copy.h>

template<typename T>
struct is_non_zero {
    __host__ __device__
    auto operator()(T x) const -> bool {
        return T != 0;
    }
};

// ... your input and output vectors here

thrust::copy_if(input.begin(), input.end(), output.begin(), is_non_zero<int>());

如果没有选择“ 推力”，则可以自己实现流压缩（有关该主题的文献很多）。这是一个有趣且相当简单的练习，同时也是更复杂的并行基元的基本构建块。

（1） 严格来说，这不是正好流传统意义上的压缩，作为流压缩传统上是稳定的算法，但你的要求不包括稳定。放宽要求可能会导致更有效的实施？

2020-07-28