Sum of following elements in array CUDA

Question

At first I'd like to notice that I've read about sum reduction, but it's not what I want to achieve.

Let's say I have array with elements:

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

Result should be:

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

In general it should be:

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

(ommiting -1 could be added later).

And I want to do this parallel in CUDA. What is the fastest way to achieve this?

Answer 1

You can use 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;
}

output

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