您所问的问题似乎是一个非常直接的流压缩问题,并且使用推力没有任何特别的问题,但是有一些曲折。为了选择要复制的行,您需要有一个流压缩算法可以使用的模板或密钥。这需要使用要复制的行列表通过搜索或选择操作来构建。
执行此操作的一个示例过程如下所示:
- 构造一个迭代器,它返回输入矩阵中任何条目的行号。推力有一个非常有用的
counting_iterator
,transform_iterator
可以结合起来做到这一点
- 执行对该行号迭代器的搜索,以查找哪些条目与要复制的行列表匹配。
thrust::binary search
可用于此。搜索产生流压缩操作的模板
- 用于使用
thrust::copy_if
模板对输入矩阵执行流压缩。
这听起来像是很多工作和中间步骤,但计数和转换迭代器实际上并没有产生任何中间设备向量。唯一需要的中间存储是模板数组,它可以是一个布尔值(所以 m*n 字节)。
代码中的完整示例:
#include <thrust/copy.h>
#include <thrust/binary_search.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/device_vector.h>
#include <cstdio>
struct div_functor : public thrust::unary_function<int,int>
{
int m;
div_functor(int _m) : m(_m) {};
__host__ __device__
int operator()(int x) const
{
return x / m;
}
};
struct is_true
{
__host__ __device__
bool operator()(bool x) { return x; }
};
int main(void)
{
// dimensions of the problem
const int m=20, n=5, l=4;
// Counting iterator for generating sequential indices
// Sample matrix containing 0...(m*n)
thrust::counting_iterator<float> indices(0.f);
thrust::device_vector<float> in_matrix(m*n);
thrust::copy(indices, indices+(m*n), in_matrix.begin());
// device vector contain rows to select
thrust::device_vector<int> select(l);
select[0] = 1;
select[1] = 4;
select[2] = 9;
select[3] = 16;
// construct device iterator supplying row numbers via a functor
typedef thrust::counting_iterator<int> counter;
typedef thrust::transform_iterator<div_functor, counter> rowIterator;
rowIterator rows_begin = thrust::make_transform_iterator(thrust::make_counting_iterator(0), div_functor(n));
rowIterator rows_end = rows_begin + (m*n);
// constructor a stencil array which indicates which entries will be copied
thrust::device_vector<bool> docopy(m*n);
thrust::binary_search(select.begin(), select.end(), rows_begin, rows_end, docopy.begin());
// use stream compaction on the matrix with the stencil array
thrust::device_vector<float> out_matrix(l*n);
thrust::copy_if(in_matrix.begin(), in_matrix.end(), docopy.begin(), out_matrix.begin(), is_true());
for(int i=0; i<(l*n); i++) {
float val = out_matrix[i];
printf("%i %f\n", i, val);
}
}
(通常的免责声明:使用风险自负)
关于我要发表的唯一评论是,copy_if
调用的谓词感觉有点多余,因为我们已经有一个可以直接使用的二进制模板,但似乎没有压缩算法的变体可以在二进制模板直接。同样,我想不出一种直接在流压缩调用中使用行列表的明智方法。很可能有一种更有效的方法可以通过推力来做到这一点,但这至少应该让你开始。
从您的评论来看,空间似乎很紧张,并且二进制搜索和模板创建的额外内存开销对于您的应用程序来说是禁止的。在这种情况下,我会遵循我在对 Roger Dahl 的回答的评论中提供的建议,并改用自定义复制内核。推力设备向量可以转换为可以直接传递给内核 ( thrust::raw_pointer_cast
) 的指针,因此它不需要干扰现有的推力代码。我建议每行使用一个线程块来复制,这允许合并读取和写入,并且应该比使用thrust::copy
每行更好。一个非常简单的实现可能看起来像这样(重用我的大部分推力示例):
#include <thrust/copy.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/device_vector.h>
#include <cstdio>
__global__
void rowcopykernel(const float *in, float *out, const int *list, const int m, const int n, const int l)
{
__shared__ const float * inrowp;
__shared__ float * outrowp;
if (threadIdx.x == 0) {
inrowp = (blockIdx.x < l) ? in + (n*list[blockIdx.x]) : 0;
outrowp = out + (n*blockIdx.x);
}
__syncthreads();
for(int i=threadIdx.x; (inrowp != 0) && (i<n); i+=blockDim.x) {
*(outrowp+i) = *(inrowp+i);
}
}
int main(void)
{
// dimensions of the problem
const int m=20, n=5, l=4;
// Sample matrix containing 0...(m*n)
thrust::counting_iterator<float> indices(0.f);
thrust::device_vector<float> in_matrix(m*n);
thrust::copy(indices, indices+(m*n), in_matrix.begin());
// device vector contain rows to select
thrust::device_vector<int> select(l);
select[0] = 1;
select[1] = 4;
select[2] = 9;
select[3] = 16;
// Output matrix
thrust::device_vector<float> out_matrix(l*n);
// raw pointer to thrust vectors
int * selp = thrust::raw_pointer_cast(&select[0]);
float * inp = thrust::raw_pointer_cast(&in_matrix[0]);
float * outp = thrust::raw_pointer_cast(&out_matrix[0]);
dim3 blockdim = dim3(128);
dim3 griddim = dim3(l);
rowcopykernel<<<griddim,blockdim>>>(inp, outp, selp, m, n, l);
for(int i=0; i<(l*n); i++) {
float val = out_matrix[i];
printf("%i %f\n", i, val);
}
}
(标准免责声明:使用风险自负)。
执行参数的选择可以做得更漂亮,但除此之外,这应该是所需的全部。如果您的行非常小,您可能希望使用每行的扭曲而不是块进行调查(因此一个块复制几行)。如果您有超过 65535 个输出行,那么您将需要使用 2D 网格,或者修改代码以使每个块执行多行。但是,与基于推力的解决方案一样,这应该可以帮助您入门。