CUDA流压缩算法

Question

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

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

Answer 1

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

如果Thrust是一个选项,你可以简单地使用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>());

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

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

Answer 2

流压缩是一个众所周知的问题，为此编写了大量代码（Thrust、Chagg 引用了两个在 CUDA 上实现流压缩的库）。

如果您有一个相对较新的支持 CUDA 的设备，支持内部函数__ballot（计算能力 >= 3.0），那么值得尝试一个比 Thrust 更快地执行流压缩的小型 CUDA过程。

这里找到代码和最小文档。 https://github.com/knotman90/cuStreamComp

它以单内核方式使用投票函数来执行压缩。

编辑：

我写了一篇文章解释这种方法的内部工作原理。 如果您有兴趣，可以在这里找到它。

Answer 3

通过这个答案，我只是想为 Davide Spataro 的方法提供更多细节。

\n\n

正如您所提到的，流压缩包括根据谓词删除集合中不需要的元素。例如，考虑一个整数数组和谓词p(x)=x>5，该数组A={6,3,2,11,4,5,3,7,5,77,94,0}被压缩为B={6,11,7,77,94}。

\n\n

流压缩方法的总体思想是将不同的计算线程分配给要压缩的数组的不同元素。每个这样的线程必须决定将其相应的元素写入输出数组，具体取决于它是否满足相关谓词。因此，流压缩的主要问题是让每个线程知道相应的元素必须写入输出数组中的哪个位置。

\n\n

copy_if[1,2] 中的方法是上述Thrust 的替代方法，由三个步骤组成：

\n\n

\n

1. 步骤1。令P为已启动线程的数量N，其中 为N>P要压缩的向量的大小。输入向量被划分为大小S等于块大小的子向量。利用块__syncthreads_count(pred)内在函数来计算块中满足谓词 pred 的线程数。作为第一步的结果，数组 的每个元素d_BlockCounts（其大小为N/P）包含相应块中满足谓词 pred 的元素数量。

\n

2. 第2步。对数组 d_BlockCounts 执行独占扫描操作。作为第二步的结果，每个线程都知道前面的块中有多少个元素写入了一个元素。因此，它知道写入其相应元素的位置，但偏移量与其自己的块相关。

\n

3. 步骤＃3。每个线程使用 warp 内部函数计算提到的偏移量，并最终写入输出数组。需要注意的是，步骤#3的执行与warp调度有关。因此，输出数组中的元素顺序不一定反映输入数组中的元素顺序。

\n

\n\n

在上述三个步骤中，第二个步骤由 CUDA Thrust\xe2\x80\x99sexclusive_scan原语执行，并且计算要求明显低于其他两个步骤。

\n\n

对于元素数组2097152，上述方法已在卡0.38ms上执行，这与CUDA Thrust\xe2\x80\x99sNVIDIA GTX 960不同。上述方法似乎更快，原因有二：\n1) 它是专门为支持 warp 内在元素的卡定制的；\n2) 该方法不保证输出顺序。1.0mscopy_if

\n\n

应该注意的是，我们还针对inkc.sourceforge.net上提供的代码测试了该方法。尽管后一种代码被安排在单个内核调用中（它不使用任何 CUDA Thrust 原语），但与三内核版本相比，它的性能并不好。

\n\n

完整代码可在此处获取，与原始 Davide Spataro 的例程相比略有优化。

\n\n

[1] M.Biller, O. Olsson, U. Assarsson, \xe2\x80\x9cEfficient stream compaction on wide SIMD many-core architectures,\xe2\x80\x9d Proc. of the Conf. on High Performance Graphics, New Orleans, LA, Aug. 01 - 03, 2009, pp. 159-166.\n[2] D.M. Hughes, I.S. Lim, M.W. Jones, A. Knoll, B. Spencer, \xe2\x80\x9cInK-Compact: in-kernel stream compaction and its application to multi-kernel data visualization on General-Purpose GPUs,\xe2\x80\x9d Computer Graphics Forum, vol. 32, n. 6, pp. 178-188, 2013.\n

\n