采用迭代器元组并生成基本类型的一元函数

Question

我有一个指向迭代器元组的 zip 迭代器。我想为 thrust::transform 提供一个仿函数，它将使用元组获取元素并生成标量输出。

我的程序不工作，我不知道为什么。

我认为这可能与以下问题有关：，但更改仿函数的模板参数并没有解决问题。

编译以下代码：

#include <iostream>
#include <thrust/transform.h>
#include <thrust/functional.h>
#include <thrust/device_vector.h>
#include <thrust/device_ptr.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/tuple.h>

typedef thrust::device_vector<double>::iterator realIter;
typedef thrust::tuple<realIter,realIter> Tup;
typedef thrust::zip_iterator<Tup> Zip;
typedef thrust::tuple<double,double> Tup2; //I tried replacing Tup with this in the functor

struct dummyOp : public thrust::unary_function<Tup, double> {
 __host__ __device__ double operator()(Tup &tup){
   double result = *thrust::get<0>(tup);
   return result;
 }
};

int main(){
  thrust::device_vector<double> A(4);
  thrust::device_vector<double> B(4);
  thrust::device_vector<double> C(4);

  A[0] =  1.; A[1] = 2.;
  A[2] =  3.; A[3] = 4.;

  B[0] =  4.;   B[1] = 3.;
  B[2] =  2.;   B[3] = 1.;

  Tup tup   = thrust::tuple<realIter,realIter>(A.begin(),B.begin());
  Zip zippy = thrust::zip_iterator<Tup>(tup);
  dummyOp f;

  // The following does not work:
  //thrust::transform(zippy, zippy + 4, C.begin(), f);

  std::cout << "A:\n";
  thrust::copy(A.begin(), A.end(), std::ostream_iterator<double>(std::cout, " "));
  std::cout << "\nB:\n";
  thrust::copy(B.begin(), B.end(), std::ostream_iterator<double>(std::cout, " "));
  std::cout << "\nC:\n";
  thrust::copy(C.begin(), C.end(), std::ostream_iterator<double>(std::cout, " "));
  std::cout << std::endl;
  std::cout <<"get<0>(zippy[0]) returns:\n" << thrust::get<0>(zippy[0]) << std::endl;
  std::cout <<"get<1>(zippy[1]) returns:\n" << thrust::get<1>(zippy[1]) << std::endl;

  return 0;
}

和运行它给出：

$ ./so2
A:
1 2 3 4
B:
4 3 2 1
C:
0 0 0 0
get<0>(zippy[0]) returns:
1
get<1>(zippy[1]) returns:
3

取消注释违规行，我们有：

$ nvcc -arch=compute_35 so2.cu -o so2
/usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/detail/internal_functional.h(322): error: function "dummyOp::operator()" cannot be called with the given argument list
            argument types are: (thrust::detail::tuple_of_iterator_references<double &, double &, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>)
            object type is: dummyOp
          detected during:
            instantiation of "thrust::detail::enable_if_non_const_reference_or_tuple_of_iterator_references<thrust::tuple_element<1, Tuple>::type>::type thrust::detail::unary_transform_functor<UnaryFunction>::operator()(Tuple) [with UnaryFunction=dummyOp, Tuple=thrust::detail::tuple_of_iterator_references<thrust::detail::tuple_of_iterator_references<double &, double &, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>, double &, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>]"
/usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/detail/function.h(60): here
            instantiation of "Result thrust::detail::wrapped_function<Function, Result>::operator()(const Argument &) const [with Function=thrust::detail::unary_transform_functor<dummyOp>, Result=void, Argument=thrust::detail::tuple_of_iterator_references<thrust::detail::tuple_of_iterator_references<thrust::device_reference<double>, thrust::device_reference<double>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>, thrust::device_reference<double>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>]"
/usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/system/cuda/detail/for_each.inl(57): here
            instantiation of "void thrust::system::cuda::detail::for_each_n_detail::for_each_kernel::operator()(thrust::system::cuda::detail::bulk_::parallel_group<thrust::system::cuda::detail::bulk_::concurrent_group<thrust::system::cuda::detail::bulk_::agent<1UL>, 0UL>, 0UL> &, Iterator, Function, Size) [with Iterator=thrust::zip_iterator<thrust::tuple<Zip, thrust::detail::normal_iterator<thrust::device_ptr<double>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, Function=thrust::detail::wrapped_function<thrust::detail::unary_transform_functor<dummyOp>, void>, Size=unsigned int]"
/usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/system/cuda/detail/bulk/detail/apply_from_tuple.hpp(71): here
            instantiation of "void thrust::system::cuda::detail::bulk_::detail::apply_from_tuple(Function, const thrust::tuple<Arg1, Arg2, Arg3, Arg4, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type> &) [with Function=thrust::system::cuda::detail::for_each_n_detail::for_each_kernel, Arg1=thrust::system::cuda::detail::bulk_::parallel_group<thrust::system::cuda::detail::bulk_::concurrent_group<thrust::system::cuda::detail::bulk_::agent<1UL>, 0UL>, 0UL> &, Arg2=thrust::zip_iterator<thrust::tuple<Zip, thrust::detail::normal_iterator<thrust::device_ptr<double>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, Arg3=thrust::detail::wrapped_function<thrust::detail::unary_transform_functor<dummyOp>, void>, Arg4=unsigned int]"
/usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/system/cuda/detail/bulk/detail/closure.hpp(50): here
            instantiation of "void thrust::system::cuda::detail::bulk_::detail::closure<Function, Tuple>::operator()() [with Function=thrust::system::cuda::detail::for_each_n_detail::for_each_kernel, Tuple=thrust::tuple<thrust::system::cuda::detail::bulk_::parallel_group<thrust::system::cuda::detail::bulk_::concurrent_group<thrust::system::cuda::detail::bulk_::agent<1UL>, 0UL>, 0UL> &, thrust::zip_iterator<thrust::tuple<Zip, thrust::detail::normal_iterator<thrust::device_ptr<double>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, thrust::detail::wrapped_function<thrust::detail::unary_transform_functor<dummyOp>, void>, unsigned int, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>]"
/usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/system/cuda/detail/bulk/detail/cuda_task.hpp(58): here
            [ 9 instantiation contexts not shown ]
            instantiation of "RandomAccessIterator thrust::system::cuda::detail::for_each(thrust::system::cuda::detail::execution_policy<DerivedPolicy> &, RandomAccessIterator, RandomAccessIterator, UnaryFunction) [with DerivedPolicy=thrust::system::cuda::detail::tag, RandomAccessIterator=thrust::zip_iterator<thrust::tuple<Zip, thrust::detail::normal_iterator<thrust::device_ptr<double>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, UnaryFunction=thrust::detail::unary_transform_functor<dummyOp>]"
/usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/detail/for_each.inl(44): here
            instantiation of "InputIterator thrust::for_each(const thrust::detail::execution_policy_base<DerivedPolicy> &, InputIterator, InputIterator, UnaryFunction) [with DerivedPolicy=thrust::system::cuda::detail::tag, InputIterator=thrust::zip_iterator<thrust::tuple<Zip, thrust::detail::normal_iterator<thrust::device_ptr<double>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, UnaryFunction=thrust::detail::unary_transform_functor<dummyOp>]"
/usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/system/detail/generic/transform.inl(57): here
            instantiation of "OutputIterator thrust::system::detail::generic::transform(thrust::execution_policy<DerivedPolicy> &, InputIterator, InputIterator, OutputIterator, UnaryFunction) [with DerivedPolicy=thrust::system::cuda::detail::tag, InputIterator=Zip, OutputIterator=thrust::detail::normal_iterator<thrust::device_ptr<double>>, UnaryFunction=dummyOp]"
/usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/detail/transform.inl(44): here
            instantiation of "OutputIterator thrust::transform(const thrust::detail::execution_policy_base<DerivedPolicy> &, InputIterator, InputIterator, OutputIterator, UnaryFunction) [with DerivedPolicy=thrust::system::cuda::detail::tag, InputIterator=Zip, OutputIterator=thrust::detail::normal_iterator<thrust::device_ptr<double>>, UnaryFunction=dummyOp]"
/usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/detail/transform.inl(142): here
            instantiation of "OutputIterator thrust::transform(InputIterator, InputIterator, OutputIterator, UnaryFunction) [with InputIterator=Zip, OutputIterator=thrust::detail::normal_iterator<thrust::device_ptr<double>>, UnaryFunction=dummyOp]"
so2.cu(36): here

更新：虽然我仍然很想帮助理解上面代码中的问题，但以下工作：

struct dummyOp {
 template <typename Tuple>
 __host__ __device__ double operator()(Tuple tup){
   double result = thrust::get<0>(tup);
   return result;
 }
};

这个想法是从这里偷来的：https://github.com/thrust/thrust/blob/master/examples/arbitrary_transformation.cu

这不构成对我自己问题的回答，因为我仍然不明白原始代码中的类型有什么问题。

另外：使参数成为显式引用（元组&）不起作用。是不是说tup是按值传递的？

Answer 1

一般来说，当 zip 迭代器作为 thrust 算法的一部分被取消引用时，它会创建一个传递给相关函子的基本类型元组（即不是迭代器或指针）。

当我们使用 "offending line":

分析您代码的编译器输出时

argument types are: (thrust::detail::tuple_of_iterator_references<double &, double &,

我们观察到，当 thrust 取消引用 zippy zip 迭代器时，它会生成一个包含 references 到 double 项的元组。我们可以使用它来告知我们自己期望的函子输入类型，即：

thrust::tuple<double &, double &>

因为这些是对基本类型的引用，所以我们没有必要在仿函数中取消引用它们（就像它们是指针或迭代器一样）来获取它们的值。

以下修改后的代码结合了这些思想并编译无误：

$ cat t4.cu
#include <iostream>
#include <thrust/transform.h>
#include <thrust/functional.h>
#include <thrust/device_vector.h>
#include <thrust/device_ptr.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/tuple.h>

typedef thrust::device_vector<double>::iterator realIter;
typedef thrust::tuple<realIter,realIter> Tup;
typedef thrust::zip_iterator<Tup> Zip;
typedef thrust::tuple<double &,double &> Tup2; //I tried replacing Tup with this in the functor

struct dummyOp : public thrust::unary_function<Tup2, double> {
 __host__ __device__ double operator()(Tup2 &tup){
   double result = thrust::get<0>(tup);
   return result;
 }
};

int main(){
  thrust::device_vector<double> A(4);
  thrust::device_vector<double> B(4);
  thrust::device_vector<double> C(4);

  A[0] =  1.; A[1] = 2.;
  A[2] =  3.; A[3] = 4.;

  B[0] =  4.;   B[1] = 3.;
  B[2] =  2.;   B[3] = 1.;

  Tup tup   = thrust::tuple<realIter,realIter>(A.begin(),B.begin());
  Zip zippy = thrust::zip_iterator<Tup>(tup);
  dummyOp f;

  // The following does not work:
  thrust::transform(zippy, zippy + 4, C.begin(), f);

  std::cout << "A:\n";
  thrust::copy(A.begin(), A.end(), std::ostream_iterator<double>(std::cout, " "));
  std::cout << "\nB:\n";
  thrust::copy(B.begin(), B.end(), std::ostream_iterator<double>(std::cout, " "));
  std::cout << "\nC:\n";
  thrust::copy(C.begin(), C.end(), std::ostream_iterator<double>(std::cout, " "));
  std::cout << std::endl;
  std::cout <<"get<0>(zippy[0]) returns:\n" << thrust::get<0>(zippy[0]) << std::endl;
  std::cout <<"get<1>(zippy[1]) returns:\n" << thrust::get<1>(zippy[1]) << std::endl;

  return 0;
}
$ nvcc -arch=sm_61 -o t4 t4.cu
$ ./t4
A:
1 2 3 4
B:
4 3 2 1
C:
1 2 3 4
get<0>(zippy[0]) returns:
1
get<1>(zippy[1]) returns:
3
$

作为替代方案，在元组类型上模板化仿函数运算符当然是可行的，因为编译器推导出必要的类型细节并实例化仿函数运算符的适当版本。

采用迭代器元组并生成基本类型的一元函数

Unary function that takes tuple of iterators and produces basic type

c++

thrust