数组 CUDA 中以下元素的总和

Question

首先我想注意到我读过 sum reduction，但这不是我想要达到的目标。

假设我有一个包含元素的数组：

array[0] = 2
array[1] = 1
array[2] = -1
array[3] = 3
array[4] = -1
array[5] = 2

结果应该是：

array[0] = array[0] = 2
array[1] = array[0] + array[1] = 2 + 1 = 3
array[2] = -1 (when -1 omitting)
array[3] = array[0] + array[1] + array[3] = 2 + 1 + 3 = 6 (ommited array[2] due to -1 value)
array[4] = -1 (omitting)
array[5] = array[0] + array[1] + array[3] + array[5] = 2 + 1 + 3 + 2 = 8

一般来说应该是：

array[0] = array[0]
array[1] = array[0] + array[1]
array[2] = array[1] + array[2]
array[3] = array[2] + array[3]
array[4] = array[3] + array[4]
array[5] = array[4] + array[5]

or just

array[n] += array[n-1] where n > 0

（省略 -1 可以稍后添加）。

我想在 CUDA 中并行执行此操作。实现这一目标的最快方法是什么？

Answer 1

你可以使用thrust::inclusive_scan:

#include <thrust/device_vector.h>
#include <thrust/scan.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/iterator/counting_iterator.h>
#include <iostream>

struct omit_negative : public thrust::unary_function<int, int>
{
  __host__ __device__
  int operator()(int value)
  {
    if (value<0)
    {
      value = 0;
    }
    return value;
  }
};

int main()
{
  int array[] = {2,1,-1,3,-1,2};
  const int array_size = sizeof(array)/sizeof(array[0]);
  thrust::device_vector<int> d_array(array, array + array_size);
  thrust::device_vector<int> d_result(array_size);

  std::cout << "input data" << std::endl;
  thrust::copy(d_array.begin(), d_array.end(), std::ostream_iterator<int>(std::cout, " "));

  thrust::inclusive_scan(thrust::make_transform_iterator(d_array.begin(), omit_negative()),
                         thrust::make_transform_iterator(d_array.end(),   omit_negative()),
                         d_result.begin());

  std::cout << std::endl << "after inclusive_scan" << std::endl;
  thrust::copy(d_result.begin(), d_result.end(), std::ostream_iterator<int>(std::cout, " "));

  using namespace thrust::placeholders;
  thrust::scatter_if(d_array.begin(),
                     d_array.end(),
                     thrust::make_counting_iterator(0),
                     d_array.begin(),
                     d_result.begin(),
                     _1<0
                    );

  std::cout << std::endl << "after scatter_if" << std::endl;
  thrust::copy(d_result.begin(), d_result.end(), std::ostream_iterator<int>(std::cout, " "));
  std::cout << std::endl;
}

输出

input data
2 1 -1 3 -1 2 
after inclusive_scan
2 3 3 6 6 8 
after scatter_if
2 3 -1 6 -1 8