上一教程： 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())));
}
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())));
}
我们的目标是创建一个迭代器，它将从矩阵的开头开始，并正确递增以访问连续的矩阵元素。对于连续的行来说，这是微不足道的，但是对于带步长的矩阵的一列呢？为此，我们需要迭代器知道矩阵的维度和步长。此信息嵌入在 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 operator()(int x) const
{
int row = x / columns;
int idx = (row * step) + (x % columns)*channels;
return idx;
}
};
step functor 获取一个索引值并返回从矩阵开头开始的适当偏移量。计数迭代器只是递增像素元素的范围。组合到 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 operator()(const unsigned int n) const
{
thrust::default_random_engine rng;
thrust::uniform_real_distribution<float> dist(a, b);
rng.discard(n);
return dist(rng);
}
};
这将获取一个整数值并输出 a 和 b 之间的值。现在我们将使用 thrust transform 将矩阵填充为 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 兼容的 begin 和 end 迭代器，指向此矩阵的通道 1
auto keyBegin = GpuMatBeginItr<int>(d_data, 1);
auto keyEnd = GpuMatEndItr<int>(d_data, 1);
// 将此矩阵的第 0 个通道的Thrust兼容的起始和结束迭代器
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);
}
首先，我们将使用流填充一个 GPU 矩阵，其中包含在 -1 和 1 之间随机生成的数据。
// 与之前的随机生成代码相同，只是现在变换是在流上执行的
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 代码。截至 7.5 版，CUDA 工具包附带的 thrust 版本中存在一个 bug，此 bug 尚未修复。此 bug 会导致代码无法在流上执行。但是，可以使用 git 存储库中的最新版本的 thrust 来修复此 bug。（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.12.0