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

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

我正在寻找一种快速减少多个等长块的方法 排列成一个大向量。 我有 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

您要求的是分段还原。这可以使用 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" ]

奇特的迭代器将合并为:

  1. 产生线性序列 0 1 2 3 ...(通过 counting_iterator)
  2. 将线性序列的每个成员除以 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,
$