cuda - 使用 CUDA Thrust 多次复制向量

Question

我正在尝试使用 CUDA Thrust 解决问题。

我有一个带有 3 的主机阵列元素。是否可以使用 Thrust 创建 384 的设备数组？其中3的元素我的主机数组中的元素重复 128次( 128 x 3 = 384 )？

一般来说，从3的数组开始元素，如何使用 Thrust 生成大小为 X 的设备数组，其中 X = Y x 3 ，即 Y是重复次数？

Answer 1

一种可能的方法:

创建适当大小的设备向量

创建 3 strided ranges , 一个用于最终输出(设备)向量中的每个元素位置 {1, 2, 3}

使用推力::填充用适当的(宿主向量)元素 {1, 2, 3}

填充 3 个跨步范围中的每一个。

此代码是对 strided range 示例的简单修改以进行演示。您可以更改 REPS定义为 128 以查看完整扩展到 384 个输出元素:

#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/iterator/permutation_iterator.h>
#include <thrust/functional.h>

#include <thrust/fill.h>
#include <thrust/device_vector.h>
#include <thrust/host_vector.h>

// for printing
#include <thrust/copy.h>
#include <ostream>


#define STRIDE 3
#define REPS  15  // change to 128 if you like
#define DSIZE (STRIDE*REPS)

// this example illustrates how to make strided access to a range of values
// examples:
//   strided_range([0, 1, 2, 3, 4, 5, 6], 1) -> [0, 1, 2, 3, 4, 5, 6]
//   strided_range([0, 1, 2, 3, 4, 5, 6], 2) -> [0, 2, 4, 6]
//   strided_range([0, 1, 2, 3, 4, 5, 6], 3) -> [0, 3, 6]
//   ...

template <typename Iterator>
class strided_range
{
    public:

    typedef typename thrust::iterator_difference<Iterator>::type difference_type;

    struct stride_functor : public thrust::unary_function<difference_type,difference_type>
    {
        difference_type stride;

        stride_functor(difference_type stride)
            : stride(stride) {}

        __host__ __device__
        difference_type operator()(const difference_type& i) const
        {
            return stride * i;
        }
    };

    typedef typename thrust::counting_iterator<difference_type>                   CountingIterator;
    typedef typename thrust::transform_iterator<stride_functor, CountingIterator> TransformIterator;
    typedef typename thrust::permutation_iterator<Iterator,TransformIterator>     PermutationIterator;

    // type of the strided_range iterator
    typedef PermutationIterator iterator;

    // construct strided_range for the range [first,last)
    strided_range(Iterator first, Iterator last, difference_type stride)
        : first(first), last(last), stride(stride) {}

    iterator begin(void) const
    {
        return PermutationIterator(first, TransformIterator(CountingIterator(0), stride_functor(stride)));
    }

    iterator end(void) const
    {
        return begin() + ((last - first) + (stride - 1)) / stride;
    }

    protected:
    Iterator first;
    Iterator last;
    difference_type stride;
};

int main(void)
{
    thrust::host_vector<int> h_data(STRIDE);
    h_data[0] = 1;
    h_data[1] = 2;
    h_data[2] = 3;

    thrust::device_vector<int> data(DSIZE);

    typedef thrust::device_vector<int>::iterator Iterator;
    strided_range<Iterator> pos1(data.begin(), data.end(), STRIDE);
    strided_range<Iterator> pos2(data.begin()+1, data.end(), STRIDE);
    strided_range<Iterator> pos3(data.begin()+2, data.end(), STRIDE);

    thrust::fill(pos1.begin(), pos1.end(), h_data[0]);
    thrust::fill(pos2.begin(), pos2.end(), h_data[1]);
    thrust::fill(pos3.begin(), pos3.end(), h_data[2]);


    // print the generated data
    std::cout << "data: ";
    thrust::copy(data.begin(), data.end(), std::ostream_iterator<int>(std::cout, " "));  std::cout << std::endl;

    return 0;
}