将 cv::cuda::GpuMat 与 thrust 配合使用

Thrust 是一个功能极其强大的库，用于各种 CUDA 加速算法。然而，thrust 设计为与向量协同工作，而不是与带步长（pitched）的矩阵协同工作。下面的教程将讨论如何把 cv::cuda::GpuMat 包装成可用于 thrust 算法的 thrust 迭代器。

本教程将向你展示如何：

将 GpuMat 包装成 thrust 迭代器
用随机数填充 GpuMat
就地对 GpuMat 的一列排序
将大于 0 的值复制到一个新的 GPU 矩阵
将 stream 与 thrust 配合使用

将 GpuMat 包装成 thrust 迭代器

以下代码将为 GpuMat 生成一个迭代器。

我们的目标是得到一个从矩阵起始处开始、并能正确递增以访问连续矩阵元素的迭代器。对于连续的一行来说这很容易，但对于带步长矩阵的一列呢？为此，我们需要迭代器感知矩阵的维度和步长（step）。这些信息嵌入在 step_functor 中。 step 仿函数（functor）接收一个索引值，并返回从矩阵起始处算起的合适偏移量。计数迭代器（counting iterator）只是在像素元素范围内递增。将它们组合进 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

首先我们需要编写一个能生成随机值的仿函数。

它接收一个整数值，并输出一个介于 a 和 b 之间的值。现在我们将用 thrust transform 用介于 0 到 10 之间的值填充我们的矩阵。

就地对 GpuMat 的一列排序

让我们用随机值和一个索引来填充矩阵元素。之后我们将对随机数和索引进行排序。

在使用 stream 的同时将大于 0 的值复制到一个新的 GPU 矩阵

在这个示例中，我们将看到 cv::cuda::Streams 如何与 thrust 配合使用。遗憾的是，这个特定示例使用的函数必须将结果返回给 CPU，因此这并非 stream 的最优用法。

首先，我们将在一个 stream 上用介于 -1 到 1 之间随机生成的数据填充一个 GPU 矩阵。

请注意 thrust::system::cuda::par.on(…) 的使用，它会创建一个在 stream 上执行 thrust 代码的执行策略（execution policy）。随 CUDA 工具包分发的 thrust 版本中存在一个 bug，截至 7.5 版本仍未修复。该 bug 会导致代码不在 stream 上执行。不过，可以使用 git 仓库中最新版本的 thrust 来修复这个 bug。（http://github.com/thrust/thrust.git）接下来，我们将使用 thrust::count_if 配合以下谓词（predicate）来确定有多少值大于 0：

我们将利用这些结果创建一个用于存储所复制值的输出缓冲区，然后使用 copy_if 配合相同的谓词来填充输出缓冲区。最后，我们将把这些值下载到一个 CPU 矩阵中以供查看。