Working with interleaved data in thrust

Question

what is the best way to work with interleaved data in thrust, say I want to add the values with interleave length equal to 3, for example:

[1, 2, 3, 4, 5, 6]

would give

[6, 15]

or deinterleaving the data, so

[1, 2, 3, 4, 5, 6, 7, 8, 9]

would give

[1, 4, 7, 2, 5, 8, 3, 6, 9]

thanks

Answer 1

There are two questions here. The first asks how to perform a structured reduction on a data set, and the second asks how to reorder a data set given a mapping.

The first problem can be solved by logically partitioning the data set into a collection of regularly-sized subsets, and then performing a reduction on each subset. In thrust, this can be done by combining reduce_by_key with a transformed counting_iterator. The idea is to "key" each datum with the index of its subset. reduce_by_key sums every contiguous datum with equal key.

The second problem can be solved by permuting the order of the data set. You can do this with a call to gather. Here, a transformed counting_iterator can communicate the mapping of indices from the original array into the permuted array. You can also fuse such a gather operation with other algorithms (such as transform) using a permutation_iterator. Check the example program for ideas on how to do so.

That said, permuting an array is costly on a GPU due to memory coalescing issues, so you should do so sparingly.

---

Here's the full program solving your two problems:

#include <thrust/device_vector.h>
#include <thrust/reduce.h>
#include <thrust/gather.h>
#include <thrust/functional.h>

struct divide_by_three
  : thrust::unary_function<unsigned int, unsigned int>
{
  __host__ __device__
  unsigned int operator()(unsigned int i)
  {
    return i / 3;
  }
};

struct deinterleave_index
  : thrust::unary_function<unsigned int, unsigned int>
{
  __host__ __device__
  unsigned int operator()(unsigned int i)
  {
    return (i/3) + 3 * (i%3);
  }
};

int main()
{
  using namespace thrust;

  device_vector<int> example_one(6);
  example_one[0] = 1; example_one[1] = 2; example_one[2] = 3;
  example_one[3] = 4; example_one[4] = 5; example_one[5] = 6;

  // the result will have size two
  device_vector<int> example_one_result(2);

  // for each datum, associate an key, which is the datum's index divided by three
  // reduce the data by key
  reduce_by_key(make_transform_iterator(make_counting_iterator(0u), divide_by_three()),
                make_transform_iterator(make_counting_iterator(6u), divide_by_three()),
                example_one.begin(),
                thrust::make_discard_iterator(),
                example_one_result.begin());

  std::cout << "example one input:  [ ";
  thrust::copy(example_one.begin(), example_one.end(), std::ostream_iterator<int>(std::cout, " "));
  std::cout << "]" << std::endl;

  std::cout << "example one result: [ ";
  thrust::copy(example_one_result.begin(), example_one_result.end(), std::ostream_iterator<int>(std::cout, " "));
  std::cout << "]" << std::endl;


  device_vector<int> example_two(9);
  example_two[0] = 1; example_two[1] = 2; example_two[2] = 3;
  example_two[3] = 4; example_two[4] = 5; example_two[5] = 6;
  example_two[6] = 7; example_two[7] = 8; example_two[8] = 9;

  // the result will be the same size
  device_vector<int> example_two_result(9);

  // gather using the mapping defined by deinterleave_index
  gather(make_transform_iterator(make_counting_iterator(0u), deinterleave_index()),
         make_transform_iterator(make_counting_iterator(9u), deinterleave_index()),
         example_two.begin(),
         example_two_result.begin());

  std::cout << "example two input:  [ ";
  thrust::copy(example_two.begin(), example_two.end(), std::ostream_iterator<int>(std::cout, " "));
  std::cout << "]" << std::endl;

  std::cout << "example two result: [ ";
  thrust::copy(example_two_result.begin(), example_two_result.end(), std::ostream_iterator<int>(std::cout, " "));
  std::cout << "]" << std::endl;

  return 0;
}

And the output:

$ nvcc test.cu -run
example one input:  [ 1 2 3 4 5 6 ]
example one result: [ 6 15 ]
example two input:  [ 1 2 3 4 5 6 7 8 9 ]
example two result: [ 1 4 7 2 5 8 3 6 9 ]