CUDA流压缩算法

我正在尝试用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))相对时间内运行.我认为可能存在比这更快的算法,因为我们不需要浪费时间排序(交换)零零元素对和非零非零元素对(不需要保持顺序).

所以我不太确定哪种方法最快,我仍然认为有更好的方法可以解决这个问题.有什么建议？

小智.. 8

您要求的是一种称为流压缩¹的经典并行算法.

如果Thrust是一个选项,你可以简单地使用thrust::copy_if.这是一个稳定的算法,它保留了所有元素的相对顺序.

粗略素描:

#include 

template
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());

如果Thrust 不是一个选项,您可以自己实现流压缩(有很多关于该主题的文献).这是一个有趣且相当简单的练习,同时也是更复杂的并行原语的基本构建块.

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

您要求的是一种称为流压缩¹的经典并行算法.

如果Thrust是一个选项,你可以简单地使用thrust::copy_if.这是一个稳定的算法,它保留了所有元素的相对顺序.

粗略素描:

#include 

template
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());

如果Thrust 不是一个选项,您可以自己实现流压缩(有很多关于该主题的文献).这是一个有趣且相当简单的练习,同时也是更复杂的并行原语的基本构建块.

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