CUDA流压缩算法

Dane Bouchie 发表于 Dev

丹恩·布奇（Dane Bouchie）

我正在尝试使用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))相对时间内运行的。我认为甚至有一种算法可以比这种算法更快，因为我们不需要浪费时间排序（交换）零零元素对和非零非零元素对（不需要保留顺序）。

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

用户名

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

如果选择“推力”，则可以简单地使用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）严格来说，这不是正好流传统意义上的压缩，作为流压缩传统上是稳定的算法，但你的要求不包括稳定。放宽要求可能会导致更有效的实施？

本文收集自互联网，转载请注明来源。

如有侵权，请联系[email protected] 删除。