8

我正在尝试使用 CUDA 构建一个并行算法，该算法采用整数数组并删除所有0's，无论是否保持顺序。

例子：

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

主机内存结果：{17, 13, 14, 0, 0, ...}

最简单的方法是使用主机及时删除0's 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常数）

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

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

4

3 回答 3

12

您要的是一种经典的并行算法，称为流压缩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 x != 0;
    }
};

// ... your input and output vectors here

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

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

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

于 2015-12-03T07:28:40.487 回答
4

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

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

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

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

编辑：

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

于 2016-01-06T17:55:52.793 回答
4

有了这个答案，我只是想为 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}。

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

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

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

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

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

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

对于一组2097152元素，与CUDA Thrust0.38ms的. 上面提到的方法似乎更快有两个原因：1）它是专门为支持扭曲内在元素的卡片量身定制的；2）该方法不保证输出顺序。NVIDIA GTX 9601.0mscopy_if

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

完整代码可在此处获得，并且与原始 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.
于 2017-02-07T08:22:48.260 回答