将数据从较小的矢量复制到较大的矢量

Copying data from a smaller vector to a larger vector

本文关键字:复制 数据      更新时间:2023-10-16

我正在使用Thrust进行GPU项目。与其试图解释我在做什么,我将提供一个简单、稍微通用的场景,它更容易解释,并可能在未来帮助其他人。

假设我有一个向量,我想修改向量的每三个元素。

我能想到的两种解决方案是:

1) 使用类似推力调用的转换,每隔三个元素修改一次,可能带有谓词或其他什么。

2) 将每三个元素复制到一个较小的向量中,调用transform,将这些元素复制回原始向量的原始点。

这两者中的任何一个都可能使用Thrust吗?

有其他方法或更好的方法来实现这一目标吗?

感谢所有建议!

这两者中的任何一个都可能使用Thrust吗?

是的,两者都有可能。

有其他方法或更好的方法来实现这一目标吗?

在某种程度上,最佳方法可能会有所不同,这取决于应用程序中此数据可能发生的其他情况。但在你所概述的范围内,我认为推力跨步射程可能是一个不错的选择,可能是最好的选择。

当然,您的第一个方法是可行的,它有一个适当定义的函子来约束行为。(例如,用constant_iterator压缩数据以提供数据"索引",并使函子条件为数据在相应"索引"上的转换)。然而,它的缺点是,我们需要启动3倍于所需数量的线程(因为只有三分之一的线程在进行任何实际的向量修改)。跨步范围方法对此进行了改进,因为每个线程都将完成修改所选向量元素的工作,并且没有"浪费"线程。

这种方法仍然存在一定程度的"低效",因为由于GPU数据加载特性,我们访问的数据是原来的3倍(使用函子/谓词方法或跨步范围方法)。您的第二种方法(将每三个元素复制到一个较小的向量)可以缓解这种低效性,但您需要支付数据复制操作的成本,这将抵消单个transform操作的任何好处。然而,如果您想在这个缩小的向量上执行许多额外的步骤,那么将数据复制到较小向量的开销/成本可能会通过一系列剩余的操作来恢复,这些操作不会支付访问3倍数据的"低效率"。

然而,跨步距离方法对于将元素从较大向量复制到较小向量,或者直接对较大向量启动transform操作,但仅修改特定元素,应该仍然有用。

这是一个工作示例,基本上是对跨步范围示例的一个微不足道的修改,它演示了两种可能的方法-第一种是复制然后转换,第二种是就地转换:

$ cat t996.cu
#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/copy.h>
#include <thrust/transform.h>
#include <iostream>
#define STRIDE 3
// 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::device_vector<int> data(8);
    data[0] = 10;
    data[1] = 20;
    data[2] = 30;
    data[3] = 40;
    data[4] = 50;
    data[5] = 60;
    data[6] = 70;
    data[7] = 80;
    // print the initial data
    std::cout << "initial data: " << std::endl;
    thrust::copy(data.begin(), data.end(), std::ostream_iterator<int>(std::cout, " "));  std::cout << std::endl;
    typedef thrust::device_vector<int>::iterator Iterator;
    // create strided_range with indices [0,3,6]
    strided_range<Iterator> strided(data.begin(), data.end(), STRIDE);
    // method 1: copy data from larger vector to smaller, then transform it:
    thrust::device_vector<int> result1(data.size()/STRIDE+1);
    thrust::copy(strided.begin(), strided.end(), result1.begin());
    thrust::transform(result1.begin(), result1.end(), result1.begin(), thrust::negate<int>());
    std::cout << "method 1 result: " << std::endl;
    thrust::copy(result1.begin(), result1.end(), std::ostream_iterator<int>(std::cout, " "));
    std::cout << std::endl;
    // method 2: transform data "in-place":
    std::cout << "method 2 result: " << std::endl;
    thrust::transform(strided.begin(), strided.end(), strided.begin(), thrust::negate<int>());
    thrust::copy(data.begin(), data.end(), std::ostream_iterator<int>(std::cout, " "));  std::cout << std::endl;
    return 0;
}
$ nvcc -o t996 t996.cu
$ ./t996
initial data:
10 20 30 40 50 60 70 80
method 1 result:
-10 -40 -70
method 2 result:
-10 20 30 -40 50 60 -70 80
$