将数据从较小的矢量复制到较大的矢量
Copying data from a smaller vector to a larger vector
我正在使用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
$
- 在 C++ AMP 数组中复制数据多少次?
- 字符串是否在其构造函数中复制数据
- 在不复制数据的情况下,将double数组转换为只有double成员的structs数组
- 有没有一种方法可以在不复制数据的情况下从string_view创建字符串流
- 如何在不复制数据的情况下将 cv::Mat 转换为 2d 标准::矢量
- C++ 在节点内存储对数据对象的引用,而不复制数据
- 是否可以在不分配内存或复制数据的情况下构造对象?
- 如何在不复制数据的情况下在平面数组和多维数组之间进行转换?
- C 2D char阵列有时会在随机位置复制数据
- 在不复制数据的情况下更改 opencv mat 中的列顺序
- 有没有办法将 EGL 绘制到 /dev/fb1 而不是 /dev/fb0,而无需在树莓派上复制数据
- 使用C 和目标C从原始图像中复制数据
- 给定 myvector.start() 和 myvector.end(),我想在不复制数据的情况下创建 myvector
- 复制数据时出现奇怪错误
- 如何在 C/C++ 中从数组末尾开始复制数据
- glTexImage2D 复制数据或仅指定数据
- 如何在C++中从void指针复制数据
- 在不复制数据的情况下将OpenCV Mat发送到MATLAB工作区的方法
- 如何使用 GPU-DMA 从 GPU-CUDA 代码来复制数据
- 在矢量到矢量中复制数据<float> <Mat> (Opencv/C++)