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

您要的是一种经典的并行算法，称为stream compaction¹.

如果推力是一个选项，您可以简单地使用 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 不是一个选项，您可以自己实施流压缩（关于该主题的文献很多）。这是一个有趣且相当简单的练习，同时也是更复杂的并行原语的基本构建块。

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

Answer 2

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

如果你有一个相对较新的 CUDA-capable 设备，它支持 __ballot 的内在功能（计算能力 >= 3.0），那么值得尝试执行流压缩的小型 CUDA 程序比 Thrust 快得多。

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

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

编辑：

我写了一篇文章来解释这种方法的内部工作原理。有兴趣的可以找here

Answer 3

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

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

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

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

步骤#1。令 P 为已启动线程的数量，N 和 N>P 是要压缩的向量的大小。输入向量被分成 sub-vectors，大小 S 等于块大小。 __syncthreads_count(pred) 块内在函数被利用，它计算块中满足谓词 pred 的线程数。作为第一步的结果，大小为 N/P 的数组 d_BlockCounts 的每个元素都包含在相应块中满足谓词 pred 的元素数。
步骤#2。对数组d_BlockCounts进行独占扫描操作。作为第二步的结果，每个线程都知道前面块中有多少元素写入了一个元素。因此，它知道写入其相应元素的位置，但对于与其自身块相关的偏移量。
步骤#3。每个线程使用 warp 内部函数计算提到的偏移量，并最终写入输出数组。应该注意的是，步骤#3 的执行与 warp 调度有关。因此，输出数组中的元素顺序不一定反映输入数组中的元素顺序。

在上述三个步骤中，第二个步骤由 CUDA Thrust 的 exclusive_scan 原语执行，并且在计算上比其他两个步骤要求低得多。

对于 2097152 元素的数组，上述方法已在 NVIDIA GTX 960 卡上的 0.38ms 中执行，与 CUDA Thrust 的 1.0ms [=14] 相比=].上述方法似乎更快，原因有二： 1）专门为支持warp内在元素的卡片量身定制； 2) 该方法不保证输出顺序。

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

完整代码可用 here，与原始 Davide Spataro 的例程相比，它略有优化。

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

CUDA流压缩算法

CUDA stream compaction algorithm

algorithm

parallel-processing

cuda

stream-compaction

编辑：