cuda - 在推力中处理交错数据

Question

在推力中处理交错数据的最佳方法是什么，比如我想添加交错长度等于 3 的值，例如：

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

会给

[6, 15]

或解交织数据，所以

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

会给

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

谢谢

Answer 1

这里有两个问题。第一个询问如何对数据集执行结构化归约，第二个询问如何在给定映射的情况下重新排序数据集。

第一个问题可以通过将数据集逻辑划分为大小规则的子集的集合来解决，然后对每个子集执行归约。总而言之，这可以通过reduce_by_key与 transform结合来完成counting_iterator。这个想法是用其子集的索引来“键入”每个数据。reduce_by_key用相等的键对每个连续的数据求和。

第二个问题可以通过排列数据集的顺序来解决。您可以通过调用 来做到这一点gather。在这里，transformedcounting_iterator可以将索引从原始数组映射到置换数组。您还可以transform使用permutation_iterator. 检查示例程序以获取有关如何执行此操作的想法。

也就是说，由于内存合并问题，在 GPU 上置换数组的成本很高，因此您应该谨慎行事。

---

这是解决您的两个问题的完整程序：

#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;
}

和输出：

$ 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 ]