如何修改 zip 迭代器的内容

Question

我有 1000 个点的 XYZ 位置。我需要用矩阵来转换它们中的每一个。从一个更简单的问题开始，我尝试将每个点乘以一个常数。我也只举三点为例。我使用 thrust::zip_iterator 打包 XYZ 并使用仿函数和变换操作来修改 XYZ。

我的代码如下，但是编译出错。仿函数 modify_tuple 编译得很好。但是当它用于转换操作时，我会遇到很多错误。 我的问题是如何修改 zip 迭代器的内容？或者如何使用推力将函子应用到 XYZ 类型的元组中的所有点。我可以使用其他推力算法吗？

#include <thrust/transform.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/execution_policy.h>

typedef thrust::tuple<float,float,float> Float3;

struct modify_tuple
{
    int _factor;
    modify_tuple(int factor) : _factor(factor) { }

    __host__ __device__ Float3 operator()(Float3&a) const
    {
        Float3 b=thrust::make_tuple(_factor*thrust::get<0>(a), _factor*thrust::get<1>(a), _factor*thrust::get<2>(a));      

        return b;
    }
};

int main(void)
{
    thrust::device_vector<float> X(3);
    thrust::device_vector<float> Y(3);
    thrust::device_vector<float> Z(3);

    X[0]=0,    X[0]=1,    X[0]=2;
    Y[0]=4,    Y[0]=5,    Y[0]=6;
    Z[0]=7,    Z[0]=8,    Z[0]=9;

    typedef thrust::device_vector<float>::iterator                     FloatIterator;
    typedef thrust::tuple<FloatIterator, FloatIterator, FloatIterator> FloatIteratorTuple;
    typedef thrust::zip_iterator<FloatIteratorTuple>                   Float3Iterator;

    Float3Iterator P_first = thrust::make_zip_iterator(make_tuple(X.begin(), Y.begin(), Z.begin()));
    Float3Iterator P_last  = thrust::make_zip_iterator(make_tuple(X.end(), Y.end(), Z.end()));

// This line gives errors
    thrust::transform(thrust::device, P_first, P_last, P_first, modify_tuple(2));  

    return 0;
}

编译时发现的错误：

/usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/detail/internal_functional.h(322): error: function "modify_tuple::operator()" cannot be called with the given argument list
            argument types are: (thrust::detail::tuple_of_iterator_references<float &, float &, float &, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>)
            object type is: modify_tuple
          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=modify_tuple, Tuple=thrust::detail::tuple_of_iterator_references<thrust::detail::tuple_of_iterator_references<float &, float &, float &, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>, thrust::detail::tuple_of_iterator_references<float &, float &, float &, 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::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<modify_tuple>, Result=void, Argument=thrust::detail::tuple_of_iterator_references<thrust::detail::tuple_of_iterator_references<thrust::device_reference<float>, thrust::device_reference<float>, thrust::device_reference<float>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>, thrust::detail::tuple_of_iterator_references<thrust::device_reference<float>, thrust::device_reference<float>, thrust::device_reference<float>, 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::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<thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, 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::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<modify_tuple>, 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<thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, 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::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<modify_tuple>, 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<thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, 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::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<modify_tuple>, 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
            [ 8 instantiation contexts not shown ]
            instantiation of "RandomAccessIterator thrust::system::cuda::detail::for_each_n(thrust::system::cuda::detail::execution_policy<DerivedPolicy> &, RandomAccessIterator, Size, UnaryFunction) [with DerivedPolicy=thrust::system::cuda::detail::par_t, RandomAccessIterator=thrust::zip_iterator<thrust::tuple<thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, 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::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, Size=signed long, UnaryFunction=thrust::detail::unary_transform_functor<modify_tuple>]" 
/usr/local/cuda/bin/../targets/x86_64-linux/include/thrust/system/cuda/detail/for_each.inl(173): here
            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::par_t, RandomAccessIterator=thrust::zip_iterator<thrust::tuple<thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, 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::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<modify_tuple>]" 
/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::par_t, InputIterator=thrust::zip_iterator<thrust::tuple<thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, 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::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<modify_tuple>]" 
/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::par_t, InputIterator=thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, OutputIterator=thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, UnaryFunction=modify_tuple]" 
/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::par_t, InputIterator=thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, OutputIterator=thrust::zip_iterator<thrust::tuple<thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::detail::normal_iterator<thrust::device_ptr<float>>, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>, UnaryFunction=modify_tuple]" 
modify_zip_iterator.cu(38): here

Answer 1

让函子类型与 thrust 传递的类型和 thrust 要求的类型相匹配有时会很棘手。这是一种使用模板让编译器决定所需的确切类型的可能方法：

$ cat t334.cu
#include <thrust/transform.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/execution_policy.h>
#include <iostream>
#include <thrust/copy.h>

struct modify_tuple
{
    int _factor;
    modify_tuple(int factor) : _factor(factor) { }
template <typename T>
    __host__ __device__ thrust::tuple<float,float,float> operator()(const T a) const
    {
        thrust::tuple<float,float,float>  res = a;
        thrust::get<0>(res) *= _factor;
        thrust::get<1>(res) *= _factor;
        thrust::get<2>(res) *= _factor;
        return res;
    }
};

int main(void)
{
    thrust::device_vector<float> X(3);
    thrust::device_vector<float> Y(3);
    thrust::device_vector<float> Z(3);
    thrust::device_vector<float> RX(3);
    thrust::device_vector<float> RY(3);
    thrust::device_vector<float> RZ(3);

    X[0]=0,    X[1]=1,    X[2]=2;
    Y[0]=4,    Y[1]=5,    Y[2]=6;
    Z[0]=7,    Z[1]=8,    Z[2]=9;

    thrust::copy_n(X.begin(), 3, std::ostream_iterator<float>(std::cout, ","));
    std::cout << std::endl;
    thrust::copy_n(Y.begin(), 3, std::ostream_iterator<float>(std::cout, ","));
    std::cout << std::endl;
    thrust::copy_n(Z.begin(), 3, std::ostream_iterator<float>(std::cout, ","));
    std::cout << std::endl;
    thrust::transform(thrust::device, thrust::make_zip_iterator(thrust::make_tuple(X.begin(), Y.begin(), Z.begin())),thrust::make_zip_iterator(thrust::make_tuple(X.end(), Y.end(), Z.end())), thrust::make_zip_iterator(thrust::make_tuple(RX.begin(), RY.begin(), RZ.begin())), modify_tuple(2));
    thrust::copy_n(X.begin(), 3, std::ostream_iterator<float>(std::cout, ","));
    std::cout << std::endl;
    thrust::copy_n(Y.begin(), 3, std::ostream_iterator<float>(std::cout, ","));
    std::cout << std::endl;
    thrust::copy_n(Z.begin(), 3, std::ostream_iterator<float>(std::cout, ","));
    std::cout << std::endl;
    thrust::copy_n(RX.begin(), 3, std::ostream_iterator<float>(std::cout, ","));
    std::cout << std::endl;
    thrust::copy_n(RY.begin(), 3, std::ostream_iterator<float>(std::cout, ","));
    std::cout << std::endl;
    thrust::copy_n(RZ.begin(), 3, std::ostream_iterator<float>(std::cout, ","));
    std::cout << std::endl;

    return 0;
}
$ nvcc -o t334 t334.cu
$ ./t334
0,1,2,
4,5,6,
7,8,9,
0,1,2,
4,5,6,
7,8,9,
0,2,4,
8,10,12,
14,16,18,
$

传递给仿函数的确切基础类型似乎是迭代器引用元组，而 return 类型是 float 值的普通元组。

请注意，可以直接在函子中修改a。通常看来 a 是按值传递给函子的，但在这种情况下，推力使用的是迭代器引用元组，因此在函子中修改 a 的副作用是它将修改传递给函子的实际推力矢量值（即源矢量）。因此，我在上面的代码中选择了一个实现来避免这种情况，但是对于您的原始代码，很明显您打算对源向量进行就地修改。因此，如果需要，函子可以简化为直接修改 a。

如何修改 zip 迭代器的内容

How to modify the contents of a zip iterator

zip

iterator

compiler-errors

thrust