CUDA中的条件复制，其中数据向量比模板长

Question

我想有条件地从向量复制数据，基于模板向量，这是N倍短。模板中的每个元素将负责数据向量中的N个元素。假设向量如下所示(N=3)

data = {1,2,3,4,5,6,7,8,9}
stencil = {1,0,1}

我希望得到的结果是：

result = {1,2,3,7,8,9}

有没有办法使用推力库中的函数来实现这一点？

我知道，这其中有：

thrust::copy_if (InputIterator1 first, InputIterator1 last, InputIterator2 stencil, OutputIterator result, Predicate pred)

但这并不允许我基于模板中的一个元素从数据向量中复制N个值。

Answer 1

就像通常的情况一样，我想有很多可能的方法来做到这一点。

我想到的方法(使用copy_if)是使用stencil向量作为thrust::permutation_iterator的一部分，该向量获取stencil向量并使用thrust::transform_iterator生成索引。在这个例子中，如果我们想象一个从0..8开始的复制索引，那么我们可以使用一个使用整数除以N (使用推力占位符)的thrust::counting_iterator计算出的"map“索引来索引到”源“(即模板)向量。复制谓词只测试模板值是否为== 1。

推力quick start guide简洁地描述了如何使用这些奇特的迭代器。

下面是一个工作示例：

$ cat t471.cu
#include <thrust/copy.h>
#include <thrust/device_vector.h>
#include <thrust/iterator/permutation_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/iterator/counting_iterator.h>
#include <iostream>

using namespace thrust::placeholders;

int main(){

  int data[] = {1,2,3,4,5,6,7,8,9};
  int stencil[] = {1,0,1};
  int ds = sizeof(data)/sizeof(data[0]);
  int ss = sizeof(stencil)/sizeof(stencil[0]);
  int N = ds/ss;  // assume this whole number divisible

  thrust::device_vector<int> d_data(data, data+ds);
  thrust::device_vector<int> d_stencil(stencil, stencil+ss);
  thrust::device_vector<int> d_result(ds);
  int rs = thrust::copy_if(d_data.begin(), d_data.end(), thrust::make_permutation_iterator(d_stencil.begin(), thrust::make_transform_iterator(thrust::counting_iterator<int>(0), _1 / N)), d_result.begin(), _1 == 1) - d_result.begin();
  thrust::copy_n(d_result.begin(), rs, std::ostream_iterator<int>(std::cout, ","));
  std::cout << std::endl;
  return 0;
}
$ nvcc -o t471 t471.cu
$ ./t471
1,2,3,7,8,9,
$

有了这里关于stencil组织的假设，我们还可以使用thrust::reduce预计算结果大小rs，并使用它来分配结果向量的大小：

$ cat t471.cu
#include <thrust/copy.h>
#include <thrust/reduce.h>
#include <thrust/device_vector.h>
#include <thrust/iterator/permutation_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/iterator/counting_iterator.h>
#include <iostream>

using namespace thrust::placeholders;

int main(){

  int data[] = {1,2,3,4,5,6,7,8,9};
  int stencil[] = {1,0,1};
  int ds = sizeof(data)/sizeof(data[0]);
  int ss = sizeof(stencil)/sizeof(stencil[0]);
  int N = ds/ss;  // assume this whole number divisible

  thrust::device_vector<int> d_data(data, data+ds);
  thrust::device_vector<int> d_stencil(stencil, stencil+ss);
  int rs = thrust::reduce(d_stencil.begin(), d_stencil.end())*N;
  thrust::device_vector<int> d_result(rs);
  thrust::copy_if(d_data.begin(), d_data.end(), thrust::make_permutation_iterator(d_stencil.begin(), thrust::make_transform_iterator(thrust::counting_iterator<int>(0), _1 / N)), d_result.begin(), _1 == 1) - d_result.begin();
  thrust::copy_n(d_result.begin(), rs, std::ostream_iterator<int>(std::cout, ","));
  std::cout << std::endl;
  return 0;
}
$ nvcc -o t471 t471.cu
$ ./t471
1,2,3,7,8,9,
$