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 库中的函数实现此目的？

我知道，有：

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 开始的复制索引，那么我们可以使用 "map" 计算的 "source" 索引索引到 "source" （即模板）向量整数除以 N（使用推力占位符）。复制谓词只是测试模板值是否 == 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 组织的假设，我们还可以 pre-compute 结果大小 rs 和 thrust::reduce，并使用它来分配结果向量大小:

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

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

Conditional copying in CUDA, where data vector is longer than stencil

cuda

thrust