使用 CUDA 减少排列在一个大向量中的多个等长块

Question

我正在寻找一种快速减少多个等长块的方法排列成一个大向量。我有 N 个子数组（连续元素）排列在一个大数组中。每个子数组都有一个固定的大小：k。所以整个数组的大小是：N*K

我在做的是调用内核N次。每次它计算子数组的缩减如下：我将遍历大向量中包含的所有子数组：

    for(i=0;i<N;i++){
       thrust::device_vector< float > Vec(subarray, subarray+k);
       float sum = thrust::reduce(Vec.begin(), Vec.end(), (float)0, thrust::plus<float>());
       printf("sum %f\n",sum);
 }

对于纯 CUDA，我会这样做（伪代码）：

 for(i=0;i<N;i++){
        reduction_kernel(subarray)

         }

您是否有其他解决方案来一次性执行连续子数组的归约？使用纯 CUDA 或 Thrust

Answer 1

您要求的是分段还原。这可以使用 thrust::reduce_by_key 推力完成除了长度为 N*K 的数据向量之外，我们还需要一个定义每个段的 "key" 向量——段不必相同大小，只要键向量像这样区分段：

data:  1 3 2 3 1 4 2 3 2 1 4 2 ...
keys:  0 0 0 1 1 1 0 0 0 3 3 3 ...
seg:   0 0 0 1 1 1 2 2 2 3 3 3 ...

只要键序列发生变化，键就会划定一个新的段（请注意，我在上面的示例中有两个单独的段，它们使用相同的键划定 - thrust 不会将这些段组合在一起，而是将它们分开处理，因为有 1 个或多个不同的中间键值）。您实际上没有这些数据，但是为了速度和效率，由于您的段长度相等，我们可以使用推力 fancy iterators.[=16= 的组合来生成必要的键序列 "on the fly" ]

奇特的迭代器将合并为：

产生线性序列 0 1 2 3 ...（通过 counting_iterator）
将线性序列的每个成员除以 K，段长度（通过 transform_iterator）。我在这里使用 thrust placeholder methodology，所以我不必为变换迭代器编写仿函数。

这将产生必要的 segment-key 序列。

这是一个有效的例子：

$ cat t1282.cu
#include <thrust/reduce.h>
#include <thrust/device_vector.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/discard_iterator.h>
#include <thrust/copy.h>
#include <thrust/execution_policy.h>
#include <iostream>

const int N = 1000;  // sequences
const int K = 100;   // length of sequence
typedef int mytype;

using namespace thrust::placeholders;

int main(){

  thrust::device_vector<mytype> data(N*K, 1);
  thrust::device_vector<mytype> sums(N);
  thrust::reduce_by_key(thrust::device, thrust::make_transform_iterator(thrust::counting_iterator<int>(0), _1/K), thrust::make_transform_iterator(thrust::counting_iterator<int>(N*K), _1/K), data.begin(), thrust::discard_iterator<int>(), sums.begin());
  // just display the first 10 results
  thrust::copy_n(sums.begin(), 10, std::ostream_iterator<mytype>(std::cout, ","));
  std::cout << std::endl;
}

$ nvcc -arch=sm_35 -o t1282 t1282.cu
$ ./t1282
100,100,100,100,100,100,100,100,100,100,
$

使用 CUDA 减少排列在一个大向量中的多个等长块

Reduce multiple blocks of equal length that are arranged in a big vector Using CUDA

cuda

nvidia

gpgpu

thrust

reduction