c++ - CUDA 推力 remove_if 与 operlapping 模板序列

Question

我试图从两个 thrust::device_vector<int> 中删除元素基于第一个 vector 的值。凭直觉，我创建了以下片段:

thrust::device_vector<float> idxToValue(COUNT_MAX);
thrust::device_vector<int> idxSorted(COUNT_MAX);
thrust::device_vector<int> groupIdxSorted(COUNT_MAX);
int count = COUNT_MAX;
float const minThreshold = MIN_THRESHOLD;

auto idxToValueSortedIter = thrust::make_permutation_iterator(
    idxToValue.begin()
    , idxSorted.begin()
    );

auto new_end = thrust::remove_if(
    thrust::make_zip_iterator(thrust::make_tuple(idxSorted.begin(), groupIdxSorted.begin()))
    , thrust::make_zip_iterator(thrust::make_tuple(idxSorted.begin() + count, groupIdxSorted.begin() + count))
    , idxToValueSortedIter 
    , thrust::placeholders::_1 >= minThreshold
    );

count = thrust::get<0>(new_end.get_iterator_tuple()) - idxSorted.begin();

不幸的是，Thrust 文档说

The range [stencil, stencil + (last - first)) shall not overlap the range [result, result + (last - first))

所以就我而言 idxToValueSortedIter ，用作模板序列，取决于 idxSorted并且实际上与结果重叠(相同的 vector )。

有没有办法不用将数据复制到临时 vector 就可以解决这个问题？

Answer 1

我认为你可以通过使用 remove_if 的非模板版本来做到这一点(没有模板，它对模板与输出序列的重叠没有这样的限制)，并通过你的模板(即你的排列迭代器)作为你的 zip_iterator 到 remove_if 的第三个成员加上一个适当的选择仿函数。这是一个有效的例子:

$ cat t572.cu
#include <iostream>
#include <thrust/device_vector.h>
#include <thrust/remove.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/iterator/permutation_iterator.h>
#include <thrust/copy.h>

#define COUNT_MAX 10
#define MIN_THRESHOLD 4.5f

struct my_functor
{
  float thresh;
  my_functor(float _thresh): thresh(_thresh) {}

  template <typename T>
  __host__ __device__
  bool operator()(T &mytuple) const {
    return thrust::get<2>(mytuple) > thresh;
  }
};

int main(){

  float h_idxToValue[COUNT_MAX] = {0.0f, 1.0f, 2.0f, 3.0f, 4.0f, 5.0f, 6.0f, 7.0f, 8.0f, 9.0f};
  int   h_idxSorted[COUNT_MAX] = {9, 8, 7, 6, 5, 4, 3, 2, 1, 0};
  int   h_groupIdxSorted[COUNT_MAX] = {20, 21, 22, 23, 24, 25, 26, 27, 28, 29};

  thrust::device_vector<float> idxToValue(h_idxToValue, h_idxToValue + COUNT_MAX);
  thrust::device_vector<int> idxSorted(h_idxSorted, h_idxSorted + COUNT_MAX);
  thrust::device_vector<int> groupIdxSorted(h_groupIdxSorted, h_groupIdxSorted + COUNT_MAX);
  int count = COUNT_MAX;
  float const minThreshold = MIN_THRESHOLD;

  auto new_end = thrust::remove_if(
    thrust::make_zip_iterator(thrust::make_tuple(idxSorted.begin(), groupIdxSorted.begin(), thrust::make_permutation_iterator(idxToValue.begin(), idxSorted.begin())))
    , thrust::make_zip_iterator(thrust::make_tuple(idxSorted.begin() + count, groupIdxSorted.begin() + count, thrust::make_permutation_iterator(idxToValue.begin(), idxSorted.begin() + count)))
    , my_functor(minThreshold)
    );

  count = thrust::get<0>(new_end.get_iterator_tuple()) - idxSorted.begin();

  std::cout << "count = " << count << std::endl;
  thrust::copy_n(groupIdxSorted.begin(), count, std::ostream_iterator<int>(std::cout, ","));
  std::cout << std::endl;
  return 0;
}

$ nvcc -arch=sm_20 -std=c++11 -o t572 t572.cu
$ ./t572
count = 5
25,26,27,28,29,
$

我们通常希望使用提供的仿函数的 remove_if 函数删除 idxToValue 值大于阈值 (4.5) 的条目。但是，由于 idxSorted 中的置换迭代器和倒序序列，我们看到高于阈值的值被保留，而其他值被删除。上面的例子是使用 CUDA 6.5 和 Fedora 20 来利用实验性的 c++11 支持。