イテレータのタプルをとり、基本型を生成する単項関数

Question

イテレータのタプルを指すzipイテレータがあります。私は、タプルを使用して要素をつかんでスカラー出力を生成するファンクタをthrust :: transformに提供したいと思います。

私のプログラムは機能せず、理由はわかりません。

私はそれが何かを持っているかもしれないと思う： CUDA thrust zip_iterator tuple transform_reduceしかし、ファンクターのテンプレートパラメータを変更することは、そのトリックをしませんでした。

#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

これは自分の質問への答えを構成するものではありません。

また、引数を明示的な参照にすることはできません（Tuple &）。それはtupに値が渡されることを意味しますか？

Answer 1

一般的に言えば、zipイテレータは推力アルゴリズムの一部として参照解除されると、問題のファンクタに渡される基本タイプのタプル（イテレータまたはポインタではない）を作成します。

私たちは、「問題のあるライン」を使用してコードからコンパイラの出力を解析する場合：

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

我々は推力がzippyジップイテレータを参照解除するとき、それは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 
$ 

：

次変更されたコードは、これらのアイデアを取り入れ、エラーなしでコンパイルファンクタ演算子の適切なバージョン。