具有操作模板序列的 CUDA 推力 remove_if

Question

我试图根据第一个向量的值从两个 thrust::device_vector<int> 中删除元素。直觉上我创建了以下片段：

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();

不幸的是，推力文档说

范围 [stencil, stencil + (last - first)) 不得与范围 [result, result + (last - first)) 重叠

所以在我的情况下，用作模板序列的idxToValueSortedIter 取决于idxSorted，实际上与结果重叠（相同的向量）。

有什么办法可以在不将数据复制到临时向量的情况下解决这个问题？

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 支持。