上一篇教程： 在 GPU 上进行相似性检查（PNSR 和 SSIM）
目标
Thrust 是一个非常强大的库，适用于各种由 cuda 加速的算法。然而，Thrust 旨在与向量而不是张量矩阵配合使用。以下教程将讨论将 cv::cuda::GpuMat 封装到可与 Thrust 算法结合使用的 Thrust 迭代器中。
本教程应向您展示如何
- 将 GpuMat 封装到 Thrust 迭代器中
- 用随机数填充 GpuMat
- 就地对 GpuMat 的一列排序
- 将大于 0 的值复制到新的 gpu 矩阵
- 在 Thrust 中使用流
将 GpuMat 封装到 Thrust 迭代器中
以下代码将生成一个 GpuMat 的迭代器
template<typename T>
thrust::permutation_iterator<thrust::device_ptr<T>, thrust::transform_iterator<step_functor<T>, thrust::counting_iterator<int>>> GpuMatBeginItr(
cv::cuda::GpuMat mat,
int channel = 0)
{
if (channel == -1)
{
channel = 0;
}
return thrust::make_permutation_iterator(thrust::device_pointer_cast(mat.
ptr<T>(0) + channel),
thrust::make_transform_iterator(thrust::make_counting_iterator(0), step_functor<T>(mat.
cols, mat.
step /
sizeof(T), mat.
channels())));
}
用于 OpenCV 原始数据类型的模板“trait”类。
定义 traits.hpp:113
带引用计数的 GPU 内存的基本存储类。
定义 cuda.hpp:106
GpuMat reshape(int cn, int rows=0) const
int channels() const
返回通道数
uchar * ptr(int y=0)
返回指向第 y 行的指针
size_t step
连续行之间的距离（以字节为单位）；包括间隙（如果有）。
定义 cuda.hpp:354
#define CV_Assert(expr)
检查运行时条件，如果条件不满足，则抛出异常。
定义 base.hpp:342
template<typename T>
thrust::permutation_iterator<thrust::device_ptr<T>, thrust::transform_iterator<step_functor<T>, thrust::counting_iterator<int>>> GpuMatEndItr(
cv::cuda::GpuMat mat,
int channel = 0)
{
if (channel == -1)
{
channel = 0;
}
return thrust::make_permutation_iterator(thrust::device_pointer_cast(mat.
ptr<T>(0) + channel),
thrust::make_transform_iterator(thrust::make_counting_iterator(mat.
rows*mat.
cols), step_functor<T>(mat.
cols, mat.
step /
sizeof(T), mat.
channels())));
}
int rows
行数和列数
定义 cuda.hpp:351
我们的目标是拥有一个从矩阵开头开始，并正确递增以访问连续矩阵元素的迭代器。对于连续行，这是很简单的，但是对于倾斜矩阵的一列，该怎么做呢？为此，我们需要让迭代器意识到矩阵的维度和步长。这些信息嵌入在 step_functor 中。
template<typename T> struct step_functor : public thrust::unary_function<int, int>
{
int columns;
int step;
int channels;
__host__ __device__ step_functor(int columns_, int step_, int channels_ = 1) : columns(columns_), step(step_), channels(channels_) { };
{
step = mat.
step /
sizeof(T);
}
__host__ __device__
int 运算符()(int x) const
{
int 行 = x / 列;
int 索引 = (行 * 步长) + (x % 列)*通道;
返回 索引;
}
};
step 函数对象接受一个索引值，并返回矩阵开头的适当偏移量。计数迭代器只针对像素元素范围进行增量。组合成 transform_iterator 后，我们得到一个迭代器，它从 0 到 M*N 进行计数，并针对倾斜的 GpuMat 内存正确进行增量。但是，这里不包括任何内存位置信息，为此我们需要 thrust::device_ptr。通过将设备指针与 transform_iterator 结合使用，我们可以将 thrust 指向我们矩阵的第一个元素，并让它按需进行步长。
用随机数填充 GpuMat
现在我们有一些用于为 thrust 制作迭代器的不错函数，让我们用它们来完成 OpenCV 无法完成的一些操作。很遗憾的是，在编写这篇文档的时候，OpenCV 没有任何 Gpu 随机数生成器。值得庆幸的是，thrust 具备这个功能，在两个函数之间进行互操作也变得容易了。示例取自 http://stackoverflow.com/questions/12614164/generating-a-random-number-vector-between-0-and-1-0-using-thrust
首先，我们需要编写一个将产生我们随机值的函数对象。
struct prg
{
float a, b;
__host__ __device__
prg(float _a = 0.f, float _b = 1.f) : a(_a), b(_b) {};
__host__ __device__
float 运算符()(const unsigned int n) const
{
thrust::default_random_engine rng;
thrust::uniform_real_distribution<float> dist(a, b);
rng.discard(n);
返回 dist(rng);
}
};
这将取一个整数，然后输出 a 和 b 之间的一个值。现在，我们将用一个 thrust 变换填充 0 到 10 之间的值到我们的矩阵中。
{
cv::cuda::GpuMat d_value(1, 100, CV_32F);
auto valueBegin = GpuMatBeginItr<float>(d_value);
auto valueEnd = GpuMatEndItr<float>(d_value);
thrust::transform(thrust::make_counting_iterator(0), thrust::make_counting_iterator(d_value.cols), valueBegin, prg(-1, 1));
cv::Mat h_value(d_value);
}
在原处对 GpuMat 的一列进行排序
让我们用随机值和一个索引填充矩阵元素。之后，我们将对随机数和索引进行排序。
{
cv::cuda::GpuMat d_data(1, 100, CV_32SC2);
// Thrust 兼容的开始和结束迭代器，针对此矩阵的通道 1
auto keyBegin = GpuMatBeginItr<int>(d_data, 1);
auto keyEnd = GpuMatEndItr<int>(d_data, 1);
// Thrust 兼容的开始和结束迭代器，针对此矩阵的通道 0
auto idxBegin = GpuMatBeginItr<int>(d_data, 0);
auto idxEnd = GpuMatEndItr<int>(d_data, 0);
// 用从 0 到 100 的数字序列填充索引通道
thrust::sequence(idxBegin, idxEnd);
// 用 0 到 10 之间的随机数组填充键通道。此处使用一个计数迭代器，将每个位置的整型值作为 prg::operator() 的输入
thrust::transform(thrust::make_counting_iterator(0), thrust::make_counting_iterator(d_data.cols), keyBegin, prg(0, 10));
// 按键通道和索引通道排序，以便键和索引保持在一起
thrust::sort_by_key(keyBegin, keyEnd, idxBegin);
cv::Mat h_idx(d_data);
}
在使用流的同时，将大于 0 的值复制到新 gpu 矩阵中
在此示例中，我们将看到如何将 cv::cuda::Streams 与 thrust 配合使用。不幸的是，这个具体示例使用那些必须将结果返回给 CPU 的函数，因此这并非对流的最佳应用。
{
cv::cuda::GpuMat d_value(1, 100, CV_32F);
auto valueBegin = GpuMatBeginItr<float>(d_value);
auto valueEnd = GpuMatEndItr<float>(d_value);
cv::cuda::Stream stream;
//! [random_gen_stream]
// 与之前的随机生成代码相同，现在的区别在于转换在流上执行
thrust::transform(thrust::system::cuda::par.on(cv::cuda::StreamAccessor::getStream(stream)), thrust::make_counting_iterator(0), thrust::make_counting_iterator(d_value.cols), valueBegin, prg(-1, 1));
//! [random_gen_stream]
// 计数我们将要复制的值的数量
int count = thrust::count_if(thrust::system::cuda::par.on(cv::cuda::StreamAccessor::getStream(stream)), valueBegin, valueEnd, pred_greater<float>(0.0));
// 分配一个复制值的目标位置
cv::cuda::GpuMat d_valueGreater(1, count, CV_32F);
// 复制满足条件的值
thrust::copy_if(thrust::system::cuda::par.on(cv::cuda::StreamAccessor::getStream(stream)), valueBegin, valueEnd, GpuMatBeginItr<float>(d_valueGreater), pred_greater<float>(0.0));
cv::Mat h_greater(d_valueGreater);
}
首先，我们会在流上生成介于 -1 和 1 之间的随机数据，并将其填充到一个 GPU 矩阵中。
// 与之前的随机生成代码相同，现在的区别在于转换在流上执行
thrust::transform(thrust::system::cuda::par.on(cv::cuda::StreamAccessor::getStream(stream)), thrust::make_counting_iterator(0), thrust::make_counting_iterator(d_value.cols), valueBegin, prg(-1, 1));
请注意 thrust::system::cuda::par.on(...) 的使用，这会创建一个执行策略，用于在流上执行 thrust 代码。cuda 工具包中分发的 thrust 版本存在缺陷，截至版本 7.5 仍未修复。此缺陷导致代码无法在流上执行。不过，可以通过使用来自 git 存储库的最新版 thrust 来修复此缺陷。（http://github.com/thrust/thrust.git）接下来，我们将利用 thrust::count_if 来确定有多少大于 0 的值，使用的谓词如下：
template<typename T> struct pred_greater
{
T value;
__host__ __device__ pred_greater(T value_) : value(value_){}
__host__ __device__ bool operator()(const T& val) const
{
return val > value;
}
};
我们将使用那些结果创建一个输出缓冲区用于存储复制的值，然后用 copy_if 和相同的谓词来填充输出缓冲区。最后，我们将下载这些值到 CPU 矩阵中以供查看。
1.9.8