OpenCV-和-CUDA-GPU-加速的计算机视觉实用指南-全-

OpenCV 和 CUDA GPU 加速的计算机视觉实用指南（全）

原文：annas-archive.org/md5/b5d2ddddf00cdfdea66355a7259934ba

译者：飞龙

协议：CC BY-NC-SA 4.0

前言

计算机视觉正在改变着众多行业，OpenCV 是计算机视觉中最广泛选择的工具，它能够在多种编程语言中工作。如今，在计算机视觉中实时处理大图像的需求日益增长，这对于仅凭 OpenCV 本身来说是难以处理的。在这种情况下，图形处理单元（GPU）和 CUDA 可以提供帮助。因此，本书提供了关于将 OpenCV 与 CUDA 集成以用于实际应用的详细概述。它从解释使用 CUDA 进行 GPU 编程开始，这对于从未使用过 GPU 的计算机视觉开发者来说是必不可少的。然后，通过一些实际示例解释了使用 GPU 和 CUDA 加速 OpenCV 的过程。当计算机视觉应用需要在现实场景中使用时，它需要部署在嵌入式开发板上。本书涵盖了在 NVIDIA Jetson Tx1 上部署 OpenCV 应用，这对于计算机视觉和深度学习应用非常受欢迎。本书的最后一部分涵盖了 PyCUDA 的概念，它可供使用 Python 与 OpenCV 一起工作的计算机视觉开发者使用。PyCUDA 是一个 Python 库，它利用 CUDA 和 GPU 的强大功能进行加速。本书为使用 OpenCV 在 C++或 Python 中加速计算机视觉应用的开发者提供了一个完整的指南，采用了一种动手实践的方法。

本书面向的对象

本书是针对那些正在使用 OpenCV 的开发者，他们现在想通过利用 GPU 处理的优势来学习如何处理更复杂图像数据。大多数计算机视觉工程师或开发者在尝试实时处理复杂图像数据时都会遇到问题。这就是使用 GPU 加速计算机视觉算法可以帮助他们开发出能够在实时处理复杂图像数据的算法的地方。大多数人认为，硬件加速只能通过 FPGA 和 ASIC 设计来实现，为此，他们需要了解硬件描述语言，如 Verilog 或 VHDL。然而，在 CUDA 发明之前，这种情况是真实的，CUDA 利用了 Nvidia GPU 的力量，可以通过使用 C++和 Python 等编程语言来加速算法。本书将帮助那些开发者通过帮助他们开发实际应用来了解这些概念。本书将帮助开发者将计算机视觉应用部署在嵌入式平台，如 NVIDIA Jetson TX1 上。

本书涵盖的内容

第一章，CUDA 简介与 CUDA 入门，介绍了 CUDA 架构以及它是如何重新定义了 GPU 的并行处理能力的。讨论了 CUDA 架构在现实场景中的应用。读者被介绍到用于 CUDA 的开发环境以及如何在所有操作系统上安装它。

第二章，使用 CUDA C 进行并行编程，教读者使用 CUDA 为 GPU 编写程序。它从一个简单的 Hello World 程序开始，然后逐步构建到 CUDA C 中的复杂示例。它还涵盖了内核的工作原理以及如何使用设备属性，并讨论了与 CUDA 编程相关的术语。

第三章，线程、同步和内存，教读者关于如何在 CUDA 程序中调用线程以及多个线程如何相互通信。它描述了当多个线程并行工作时如何进行同步。它还详细描述了常量内存和纹理内存。

第四章，CUDA 的高级概念，涵盖了 CUDA 流和 CUDA 事件等高级概念。它描述了如何使用 CUDA 加速排序算法，并探讨了使用 CUDA 加速简单的图像处理函数。

第五章，使用 CUDA 支持的 OpenCV 入门，描述了在所有操作系统上安装具有 CUDA 支持的 OpenCV 库。它解释了如何使用一个简单的程序来测试这个安装。本章还比较了带有和没有 CUDA 支持执行图像处理程序的性能。

第六章，使用 OpenCV 和 CUDA 进行基本计算机视觉操作，教读者如何使用 OpenCV 编写基本的计算机视觉操作，例如图像的像素级操作、过滤和形态学操作。

第七章，使用 OpenCV 和 CUDA 进行目标检测和跟踪，探讨了使用 OpenCV 和 CUDA 加速一些实际计算机视觉应用的步骤。它描述了用于目标检测的特征检测和描述算法。本章还涵盖了使用 Haar 级联和视频分析技术（如背景减法进行目标跟踪）的加速人脸检测。

第八章，Jetson TX1 开发板简介和 Jetson TX1 上安装 OpenCV，介绍了 Jetson TX1 嵌入式平台及其如何用于加速和部署计算机视觉应用。它描述了使用 Jetpack 在 Jetson TX1 上为 Tegra 安装 OpenCV 的过程。

第九章，在 Jetson TX1 上部署计算机视觉应用，涵盖了在 Jetson Tx1 上部署计算机视觉应用。它教读者如何构建不同的计算机视觉应用以及如何将摄像头与 Jetson Tx1 接口用于视频处理应用。

第十章，开始使用 PyCUDA，介绍了 PyCUDA，这是一个用于 GPU 加速的 Python 库。它描述了在所有操作系统上的安装过程。

第十一章，使用 PyCUDA 进行工作，教读者如何使用 PyCUDA 编写程序。它详细描述了从主机到设备的数据传输和内核执行的概念。它涵盖了如何在 PyCUDA 中处理数组以及开发复杂算法。

第十二章，使用 PyCUDA 开发基本计算机视觉应用，探讨了使用 PyCUDA 开发和加速基本计算机视觉应用。它以颜色空间转换操作、直方图计算和不同的算术运算为例，描述了计算机视觉应用。

要充分利用本书

本书涵盖的示例可以在 Windows、Linux 和 macOS 上运行。所有安装说明都在书中涵盖。预期读者对计算机视觉概念和 C++、Python 等编程语言有深入理解。建议读者拥有 Nvidia GPU 硬件来执行书中涵盖的示例。

下载示例代码文件

您可以从www.packt.com的账户下载本书的示例代码文件。如果您在其他地方购买了本书，您可以访问www.packt.com/support并注册，以便将文件直接通过电子邮件发送给您。

您可以通过以下步骤下载代码文件：

在www.packt.com登录或注册。
选择“支持”选项卡。
点击“代码下载与勘误”。
在搜索框中输入书籍名称，并遵循屏幕上的说明。

文件下载后，请确保您使用最新版本的以下软件解压或提取文件夹：

WinRAR/7-Zip for Windows
Zipeg/iZip/UnRarX for Mac
7-Zip/PeaZip for Linux

该书的代码包也托管在 GitHub 上，网址为github.com/PacktPublishing/Hands-On-GPU-Accelerated-Computer-Vision-with-OpenCV-and-CUDA。如果代码有更新，它将在现有的 GitHub 仓库中更新。

我们还有其他来自我们丰富图书和视频目录的代码包可供在github.com/PacktPublishing/上获取。查看它们！

下载彩色图像

我们还提供了一份包含本书中使用的截图/图表彩色图像的 PDF 文件。您可以从这里下载：www.packtpub.com/sites/default/files/downloads/978-1-78934-829-3_ColorImages.pdf。

代码实战

访问以下链接查看代码运行的视频：

bit.ly/2PZOYcH

使用的约定

本书中使用了多种文本约定。

CodeInText：表示文本中的代码单词、数据库表名、文件夹名、文件名、文件扩展名、路径名、虚拟 URL、用户输入和 Twitter 昵称。以下是一个示例：“将下载的WebStorm-10*.dmg磁盘映像文件作为系统中的另一个磁盘挂载。”

代码块设置如下：

html, body, #map {
 height: 100%; 
 margin: 0;
 padding: 0
}

当我们希望您注意代码块中的特定部分时，相关的行或项目将以粗体显示：

[default]
exten => s,1,Dial(Zap/1|30)
exten => s,2,Voicemail(u100)
exten => s,102,Voicemail(b100)
exten => i,1,Voicemail(s0)

任何命令行输入或输出都按以下方式编写：

$ mkdir css
$ cd css

粗体：表示新术语、重要单词或屏幕上看到的单词。例如，菜单或对话框中的单词在文本中显示如下。以下是一个示例：“从管理面板中选择系统信息。”

警告或重要提示看起来像这样。

技巧和窍门看起来像这样。

联系我们

我们始终欢迎读者的反馈。

一般反馈：如果您对本书的任何方面有疑问，请在邮件主题中提及书名，并通过customercare@packtpub.com发送电子邮件给我们。

勘误表：尽管我们已经尽一切努力确保内容的准确性，但错误仍然可能发生。如果您在这本书中发现了错误，如果您能向我们报告，我们将不胜感激。请访问www.packt.com/submit-errata，选择您的书籍，点击勘误提交表链接，并输入详细信息。

盗版：如果您在互联网上以任何形式遇到我们作品的非法副本，如果您能提供位置地址或网站名称，我们将不胜感激。请通过copyright@packt.com与我们联系，并提供材料的链接。

如果您有兴趣成为作者：如果您在某个领域有专业知识，并且您有兴趣撰写或为书籍做出贡献，请访问 authors.packtpub.com.

请留下评论。一旦您阅读并使用了这本书，为何不在您购买它的网站上留下评论呢？潜在读者可以查看并使用您的客观意见来做出购买决定，我们 Packt 可以了解您对我们产品的看法，我们的作者也可以看到他们对书籍的反馈。谢谢！

如需了解 Packt 的更多信息，请访问 packt.com.

第一章：介绍 CUDA 并开始使用 CUDA

本章为您简要介绍了 CUDA 架构及其如何重新定义了 GPU 的并行处理能力。本章将演示 CUDA 架构在实际场景中的应用。对于想要通过使用通用 GPU 和 CUDA 来加速其应用程序的软件开发人员，本章将作为入门指南。本章描述了用于 CUDA 应用程序开发的开发环境，以及如何在所有操作系统上安装 CUDA 工具包。它涵盖了如何使用 CUDA C 开发基本代码，并在 Windows 和 Ubuntu 操作系统上执行。

本章将涵盖以下主题：

介绍 CUDA
CUDA 的应用
CUDA 开发环境
在 Windows、Linux 和 macOS 上安装 CUDA 工具包
开发简单代码，使用 CUDA C

技术要求

本章要求您熟悉基本的 C 或 C++ 编程语言。本章中使用的所有代码都可以从以下 GitHub 链接下载：github.com/bhaumik2450/Hands-On-GPU-Accelerated-Computer-Vision-with-OpenCV-and-CUDA/Chapter1. 代码可以在任何操作系统上执行，尽管它仅在 Windows 10 和 Ubuntu 16.04 上进行了测试。

查看以下视频，了解代码的实际应用：

bit.ly/2PTQMUk

介绍 CUDA

统一计算设备架构（CUDA）是由 NVIDIA 开发的一个非常流行的并行计算平台和编程模型。它仅支持 NVIDIA GPU。OpenCL 用于为其他类型的 GPU（如 AMD 和 Intel）编写并行代码，但它比 CUDA 复杂。CUDA 允许通过简单的编程 API 在 图形处理单元（GPUs）上创建大规模并行应用程序。使用 C 和 C++ 的软件开发人员可以通过使用 CUDA C 或 C++ 来加速他们的软件应用程序，并利用 GPU 的强大功能。用 CUDA 编写的程序与简单的 C 或 C++ 程序类似，只是增加了用于利用 GPU 并行性的关键字。CUDA 允许程序员指定 CUDA 代码的哪一部分将在 CPU 上执行，哪一部分将在 GPU 上执行。

下一节将详细描述并行计算的需求以及 CUDA 架构如何利用 GPU 的强大功能。

并行处理

近年来，消费者对单一手持设备的功能需求越来越多。因此，需要在小面积上封装越来越多的晶体管，以便快速工作并消耗最小的功率。我们需要一个高速处理器，它可以在高时钟速度、小面积和最小功耗的情况下执行多个任务。在几十年的时间里，晶体管尺寸逐渐减小，使得在单个芯片上可以封装越来越多的晶体管成为可能。这导致了时钟速度的持续提升。然而，在过去的几年里，这种情况发生了变化，时钟速度大致保持不变。那么，这是为什么？晶体管停止变得更小了吗？答案是否定的。时钟速度保持不变的主要原因是在高时钟速率下的高功耗。在小型区域内紧密排列并高速工作的小型晶体管会消耗大量功率，因此很难保持处理器的冷却。随着时钟速度在开发方面趋于饱和，我们需要一种新的计算范式来提高处理器的性能。让我们通过一个小型的现实生活例子来理解这个概念。

假设你被告知在很短的时间内挖一个非常大的洞。你有以下三种选择来按时完成这项工作：

你可以挖得更快。
你可以购买一把更好的铲子。
你可以雇佣更多的挖掘工，他们可以帮助你完成工作。

如果我们能在本例与计算范式之间建立类比，那么第一种选择类似于拥有更快的时钟。第二种选择类似于拥有更多晶体管，每个时钟周期可以完成更多的工作。但是，正如我们在上一段中讨论的，功率限制对这些两个步骤都施加了限制。第三种选择类似于拥有许多小型且简单的处理器，它们可以并行执行任务。GPU 遵循这种计算范式。它不是拥有一个能够执行复杂任务的大而强大的处理器，而是拥有许多小型且简单的处理器，它们可以并行完成工作。GPU 架构的细节将在下一节中解释。

介绍 GPU 架构和 CUDA

GeForce 256 是 NVIDIA 在 1999 年开发的第一个 GPU。最初，GPU 仅用于在显示器上渲染高端图形。它们仅用于像素计算。后来，人们意识到如果 GPU 能够进行像素计算，那么它们也能够进行其他数学计算。如今，GPU 被用于许多除了渲染图形之外的应用。这类 GPU 被称为通用型 GPU（GPGPU）。

可能接下来出现在你脑海中的问题是，CPU 和 GPU 的硬件架构差异，这使它们能够执行并行计算。CPU 拥有复杂的控制硬件和较少的数据计算硬件。复杂的控制硬件使 CPU 在性能上具有灵活性，并提供了简单的编程接口，但从功耗角度来看是昂贵的。另一方面，GPU 拥有简单的控制硬件和更多用于数据计算硬件，这使得它能够进行并行计算。这种结构使其更加节能。缺点是它有一个更加限制性的编程模型。在 GPU 计算的早期阶段，图形 API，如 OpenGL 和 DirectX，是唯一与 GPU 交互的方式。这对不熟悉 OpenGL 或 DirectX 的普通程序员来说是一个复杂的任务。这导致了 CUDA 编程架构的发展，它提供了一种简单高效的方式与 GPU 交互。关于 CUDA 架构的更多细节将在下一节中给出。

通常，任何硬件架构的性能都是通过延迟和吞吐量来衡量的。延迟是完成给定任务所需的时间，而 吞吐量 是在给定时间内完成的任务量。这些概念并不矛盾。大多数情况下，提高一个会提高另一个。从某种意义上说，大多数硬件架构都是为了提高延迟或吞吐量而设计的。例如，假设你正在邮局排队。你的目标是尽可能快地完成工作，所以你想要提高延迟，而坐在邮局窗口的员工则希望每天看到越来越多的客户。因此，员工的目标是增加吞吐量。提高一个将导致另一个的提高，在这种情况下，但双方看待这种提高的方式是不同的。

同样，普通的顺序 CPU 是为了优化延迟而设计的，而 GPU 是为了优化吞吐量而设计的。CPU 是为了以最短的时间执行所有指令，而 GPU 是为了在给定时间内执行更多的指令。这种 GPU 的设计概念使它们在图像处理和计算机视觉应用中非常有用，这是我们在这本书中要针对的应用，因为我们不介意单个像素的处理延迟。我们想要的更多是，在给定时间内处理更多的像素，这可以在 GPU 上完成。

因此，总结来说，如果我们想在相同的时钟速度和功耗下提高计算性能，就需要并行计算。GPU 通过拥有大量并行工作的简单计算单元来提供这种能力。现在，为了与 GPU 交互并利用其并行计算能力，我们需要一个简单的并行编程架构，这正是 CUDA 提供的。

CUDA 架构

本节涵盖了在 GPU 架构中进行的硬件修改以及使用 CUDA 开发的软件程序的一般结构。我们目前不会讨论 CUDA 程序的语法，但我们将介绍编写代码的步骤。本节还将涵盖一些将在整本书中使用的术语。

CUDA 架构包括专为在 GPU 上进行通用计算而设计的几个新组件，这些组件在早期架构中并不存在。它包括统一的舍入流水线，该流水线允许 GPU 芯片上所有的算术 逻辑单元（ALUs）由单个 CUDA 程序进行调度。这些 ALU 也被设计为符合 IEEE 浮点单精度和双精度标准，以便在通用应用程序中使用。指令集也针对通用计算进行了定制，而不是针对像素计算。它还允许对内存进行任意的读写访问。这些特性使得 CUDA GPU 架构在通用应用程序中非常有用。

所有 GPU 都有许多称为核心的并行处理单元。在硬件方面，这些核心被分为流处理器和 流多处理器（SMs）。GPU 有一个由这些流多处理器组成的网格。在软件方面，CUDA 程序作为一系列并行运行的多个线程执行。每个线程在不同的核心上执行。GPU 可以看作是许多块的组合，每个块可以执行许多线程。每个块绑定到 GPU 上的不同 SM。块与 SM 之间的映射方式对 CUDA 程序员来说是未知的，但这是由调度器知道并完成的。来自同一块的所有线程可以相互通信。GPU 有一个处理线程之间通信的分层内存结构，包括一个块内和多个块之间的通信。这将在接下来的章节中详细介绍。

作为一名程序员，你可能想知道 CUDA 中的编程模型是什么，代码将如何理解它应该在 CPU 还是 GPU 上执行。对于这本书，我们将假设我们有一个由 CPU 和 GPU 组成的计算平台。我们将把 CPU 及其内存称为主机，把 GPU 及其内存称为设备。CUDA 代码包含主机和设备的代码。主机代码由常规的 C 或 C++ 编译器在 CPU 上编译，设备代码由 GPU 编译器在 GPU 上编译。主机代码通过所谓的 内核调用 来调用设备代码。它将在设备上并行启动许多线程。要启动的线程数量将由程序员提供。

现在，你可能会问这种设备代码与普通 C 代码有何不同。答案是，它与普通的顺序 C 代码相似。只是这种代码是在更多的核心上并行执行的。然而，为了使此代码工作，它需要在设备内存上的数据。因此，在启动线程之前，主机将数据从主机内存复制到设备内存。线程在设备内存上的数据上工作，并将结果存储在设备内存中。最后，这些数据被复制回主机内存以进行进一步处理。总之，开发 CUDA C 程序的基本步骤如下：

在主机和设备内存中为数据分配内存。
将数据从主机内存复制到设备内存。
通过指定并行度来启动内核。
在所有线程完成后，将数据从设备内存复制回主机内存。
释放主机和设备上使用的所有内存。

CUDA 应用

在过去十年中，CUDA 经历了前所未有的增长。它被用于各种领域的大量应用中。它已经改变了多个领域的研究。在本节中，我们将探讨一些这些领域以及 CUDA 如何加速每个领域的增长：

计算机视觉应用: 计算机视觉和图像处理算法计算密集。随着越来越多的摄像头以高清格式捕捉图像，需要实时处理这些大图像。通过这些算法的 CUDA 加速，图像分割、目标检测和分类等应用可以实现每秒超过 30 帧的实时帧率性能。CUDA 和 GPU 允许更快地训练深度神经网络和其他深度学习算法；这已经改变了计算机视觉的研究。NVIDIA 正在开发多个硬件平台，如 Jetson TX1、Jetson TX2 和 Jetson TK1，这些平台可以加速计算机视觉应用。NVIDIA 驱动平台也是为自动驾驶应用而设计的平台之一。
医学成像: 医学成像领域正在广泛使用 GPU 和 CUDA 进行 MRI 图像和计算机断层扫描（CT）图像的重建和处理。这极大地缩短了这些图像的处理时间。如今，有几种设备配备了 GPU，并且有几个库可用于使用 CUDA 加速处理这些图像。
金融计算: 所有金融公司都需要在较低的成本下进行更好的数据分析，这将有助于做出明智的决策。它包括复杂的风险评估和初始及终身保证金计算，这些必须在实时完成。GPU 帮助金融公司实时进行这些类型的分析，而不会增加太多的开销成本。
生命科学、生物信息学和计算化学：模拟 DNA 基因、测序和蛋白质对接是计算密集型任务，需要高计算资源。GPU 有助于这种分析和模拟。GPU 可以比普通 CPU 快五倍以上运行常见的分子动力学、量子化学和蛋白质对接应用程序。
气象研究和预报：与 CPU 相比，几个天气预报应用程序、海洋模拟技术和海啸预测技术利用 GPU 和 CUDA 进行更快的计算和模拟。
电子设计自动化（EDA）：由于 VLSI 技术和半导体制造工艺的日益复杂，EDA 工具的性能落后于这种技术进步。这导致模拟不完整和遗漏功能错误。因此，EDA 行业一直在寻求更快的模拟解决方案。GPU 和 CUDA 加速正在帮助这个行业加快计算密集型 EDA 模拟，包括功能模拟、布局和布线、信号完整性与电磁学、SPICE 电路模拟等。
政府和军事：GPU 和 CUDA 加速也广泛应用于政府和军队。航空航天、国防和情报行业正在利用 CUDA 加速将大量数据转换为可操作信息。

CUDA 开发环境

要开始使用 CUDA 开发应用程序，您需要为其设置开发环境。设置 CUDA 开发环境有一些先决条件。这些包括以下内容：

支持 CUDA 的 GPU
NVIDIA 显卡驱动程序
标准 C 编译器
CUDA 开发套件

如何检查这些先决条件并安装它们将在以下子节中讨论。

支持 CUDA 的 GPU

如前所述，CUDA 架构仅支持 NVIDIA GPU。它不支持 AMD 和 Intel 等其他 GPU。过去十年中几乎所有的 NVIDIA GPU 都支持 CUDA 架构，可以用于开发和执行 CUDA 应用程序。CUDA 支持的 GPU 详细列表可以在 NVIDIA 网站上找到：developer.nvidia.com/cuda-gpus。如果您能在列表中找到您的 GPU，您将能够在您的 PC 上运行 CUDA 应用程序。

如果您不知道您的 PC 上安装了哪种 GPU，可以通过以下步骤找到它：

在 Windows 上：
1. 在开始菜单中，键入设备管理器并按Enter。
2. 在设备管理器中展开显示适配器。在那里，您将找到您的 NVIDIA GPU 的名称。
在 Linux 上：
1. 打开终端。
2. 运行sudo lshw -C video。

这将列出有关您的显卡的信息，通常包括其制造商和型号。

在 macOS 上：
1. 前往苹果菜单 | 关于本机 | 更多信息。
2. 在内容列表下选择“图形/显示”。在那里，你可以找到你的 NVIDIA GPU 名称。

如果你有一个启用了 CUDA 的 GPU，那么你可以继续下一步。

NVIDIA 显卡驱动程序

如果你想要与 NVIDIA GPU 硬件进行通信，那么你需要为其安装系统软件。NVIDIA 提供了一个用于与 GPU 硬件通信的设备驱动程序。如果 NVIDIA 显卡安装正确，那么这些驱动程序会自动与你的 PC 一起安装。然而，定期从 NVIDIA 网站检查驱动程序更新是一个好习惯：www.nvidia.in/Download/index.aspx?lang=en-in。你可以通过此链接选择你的显卡和操作系统来下载驱动程序。

标准 C 编译器

每当你运行 CUDA 应用程序时，它将需要两个编译器：一个用于 GPU 代码，另一个用于 CPU 代码。GPU 代码的编译器将随 CUDA 工具包的安装一起提供，这将在下一节中讨论。你还需要安装一个标准的 C 编译器来执行 CPU 代码。根据操作系统，有不同的 C 编译器：

在 Windows 上：对于所有 Microsoft Windows 版本，建议使用 Microsoft Visual Studio C 编译器。它包含在 Microsoft Visual Studio 中，并且可以从其官方网站下载：www.visualstudio.com/downloads/。

商业应用的精简版需要购买，但在非商业应用中你可以免费使用社区版。为了运行 CUDA 应用程序，请安装带有 Microsoft Visual Studio C 编译器的 Microsoft Visual Studio。不同的 CUDA 版本支持不同的 Visual Studio 版本，因此你可以参考 NVIDIA CUDA 网站以了解 Visual Studio 版本支持情况。

在 Linux 上：大多数 Linux 发行版都自带标准的 GNU C 编译器（GCC），因此它可以用来编译 CUDA 应用程序的 CPU 代码。
在 Mac 上：在 Mac 操作系统上，你可以通过下载和安装 macOS 的 Xcode 来安装 GCC 编译器。它是免费提供的，可以从苹果的网站下载：

developer.apple.com/xcode/

CUDA 开发套件

CUDA 需要一个 GPU 编译器来编译 GPU 代码。这个编译器包含在 CUDA 开发工具包中。如果你有一个带有最新驱动程序更新的 NVIDIA GPU，并且已经为你的操作系统安装了标准的 C 编译器，那么你可以继续到最后一步，安装 CUDA 开发工具包。下一节将讨论安装 CUDA 工具包的逐步指南。

在所有操作系统上安装 CUDA 工具包

本节涵盖了如何在所有支持的平台上安装 CUDA 的说明。它还描述了验证安装的步骤。在安装 CUDA 时，您可以选择网络安装程序或离线本地安装程序。网络安装程序具有较小的初始下载大小，但在安装过程中需要互联网连接。本地离线安装程序具有较大的初始下载大小。本书中讨论的步骤适用于本地安装。可以从以下链接下载适用于 Windows、Linux 和 macOS 的 CUDA 工具包，包括 32 位和 64 位架构：developer.nvidia.com/cuda-downloads。

下载安装程序后，根据您的特定操作系统参考以下步骤。步骤中使用 CUDAx.x 作为表示法，其中 x.x 表示您已下载的 CUDA 版本。

Windows

本节涵盖了在 Windows 上安装 CUDA 的步骤，如下所示：

双击安装程序。它将要求您选择临时安装文件将被提取的文件夹。选择您选择的文件夹。建议保持默认设置。
然后，安装程序将检查系统兼容性。如果您的系统兼容，您可以按照屏幕提示安装 CUDA。您可以选择快速安装（默认）或自定义安装。自定义安装允许您选择要安装的 CUDA 功能。建议选择快速默认安装。
安装程序还将安装 CUDA 示例程序和 CUDA Visual Studio 集成。

在运行此安装程序之前，请确保您已安装 Visual Studio。

为了确认安装成功，以下方面应得到保证：

如果您选择了默认安装路径，所有 CUDA 示例都将位于 C:\ProgramData\NVIDIA Corporation\CUDA Samples\vx.x。
要检查安装，您可以运行任何项目。
我们使用位于 C:\ProgramData\NVIDIA Corporation\CUDA Samples\vx.x\1_Utilities\deviceQuery 的设备查询项目。
双击您 Visual Studio 版本的 *.sln 文件。它将在 Visual Studio 中打开此项目。
然后，您可以在 Visual Studio 中点击本地 Windows 调试器。如果构建成功并显示以下输出，则表示安装完成：

Linux

本节涵盖了在 Linux 发行版上安装 CUDA 的步骤。在本节中，使用特定于发行版的软件包或使用 apt-get 命令（仅适用于 Ubuntu）讨论了在流行的 Linux 发行版 Ubuntu 中的 CUDA 安装。

使用从 CUDA 网站下载的 *.deb 安装程序安装 CUDA 的步骤如下：

打开终端并运行 dpkg 命令，该命令用于在基于 Debian 的系统中安装软件包：

sudo dpkg -i cuda-repo-<distro>_<version>_<architecture>.deb

使用以下命令安装 CUDA 公共 GPG 密钥：

sudo apt-key add /var/cuda-repo-<version>/7fa2af80.pub

然后，使用以下命令更新 apt 仓库缓存：

sudo apt-get update

然后，您可以使用以下命令安装 CUDA：

sudo apt-get install cuda

使用以下命令在 PATH 环境变量中包含 CUDA 安装路径：

如果您没有在默认位置安装 CUDA，您需要更改路径以指向您的安装位置。

 export PATH=/usr/local/cuda-x.x/bin${PATH:+:${PATH}}

设置 LD_LIBRARY_PATH 环境变量：

export LD_LIBRARY_PATH=/usr/local/cuda-x.x/lib64\
${LD_LIBRARY_PATH:+:${LD_LIBRARY_PATH}}

您还可以通过使用 Ubuntu OS 中的 apt-get 软件包管理器来安装 CUDA 工具包。您可以在终端中运行以下命令：

sudo apt-get install nvidia-cuda-toolkit

要检查 CUDA GPU 编译器是否已安装，您可以从终端运行 nvcc -V 命令。它调用 GCC 编译器来编译 C 代码，以及 NVIDIA PTX 编译器来编译 CUDA 代码。

您可以使用以下命令安装 NVIDIA Nsight Eclipse 插件，它将为执行 CUDA 程序提供 GUI 集成开发环境：

sudo apt install nvidia-nsight

安装完成后，您可以在 ~/NVIDIA_CUDA-x.x_Samples 位置运行 deviceQuery 项目。如果 CUDA 工具包已正确安装和配置，deviceQuery 的输出应类似于以下内容：

Mac

本节介绍了在 macOS 上安装 CUDA 的步骤。需要从 CUDA 网站下载的 *.dmg 安装程序。下载安装程序后的安装步骤如下：

启动安装程序，按照屏幕上的提示完成安装。它将安装所有先决条件、CUDA、工具包和 CUDA 示例。
然后，您需要使用以下命令设置环境变量，以指向 CUDA 安装位置：

如果您没有在默认位置安装 CUDA，您需要更改路径以指向您的安装位置。

 export PATH=/Developer/NVIDIA/CUDA-x.x/bin${PATH:+:${PATH}}
 export DYLD_LIBRARY_PATH=/Developer/NVIDIA/CUDA-x.x/lib\
 ${DYLD_LIBRARY_PATH:+:${DYLD_LIBRARY_PATH}}

运行脚本：cuda-install-samples-x.x.sh。它将以写权限安装 CUDA 示例。
完成后，您可以去 bin/x86_64/darwin/release 并运行 deviceQuery 项目。如果 CUDA 工具包已正确安装和配置，它将显示您的 GPU 的设备属性。

CUDA C 中的基本程序

在本节中，我们将通过编写一个非常基础的 CUDA C 程序来开始学习 CUDA 编程。我们将从编写一个 Hello, CUDA! 程序开始，并执行它。在深入代码细节之前，您应该记住的是，主机代码由标准 C 编译器编译，而设备代码由 NVIDIA GPU 编译器执行。一个 NVIDIA 工具将主机代码传递给标准 C 编译器，例如 Windows 的 Visual Studio 和 Ubuntu 的 GCC 编译器，并使用 macOS 来执行。还重要的是要注意，GPU 编译器可以在没有任何设备代码的情况下运行 CUDA 代码。所有 CUDA 代码都必须以 *.cu 扩展名保存。

以下为 Hello, CUDA! 的代码：

#include <iostream>
 __global__ void myfirstkernel(void) {
 }
int main(void) {
  myfirstkernel << <1, 1 >> >();
  printf("Hello, CUDA!\n");
  return 0;
}

如果你仔细查看代码，它看起来会非常类似于为 CPU 执行编写的简单 Hello, CUDA! C 程序。这段代码的功能也是相似的。它只是在终端或命令行上打印 Hello, CUDA!。所以，应该出现在你脑海中的两个问题是：这段代码有什么不同，CUDA C 在这段代码中扮演什么角色？这些问题的答案可以通过仔细查看代码来给出。与简单 C 编写的代码相比，它有两个主要的不同点：

带有 __global__ 前缀的空函数 myfirstkernel
使用 << <1,1> >> 调用 myfirstkernel 函数

__global__ 是 CUDA C 添加到标准 C 中的一个限定符。它告诉编译器，此限定符之后的功能定义应该被编译在设备上运行，而不是在主机上。因此，在前面的代码中，myfirstkernel 将在设备上运行而不是在主机上，尽管在这个代码中它是空的。

那么，主函数将在哪里运行？NVCC 编译器会将此函数传递给宿主 C 编译器，因为它没有被 global 关键字装饰，因此 main 函数将在宿主上运行。

代码中的第二个不同点是调用空的 myfirstkernel 函数，并带有一些尖括号和数字值。这是从宿主代码调用设备代码的 CUDA C 技巧。这被称为 kernel 调用。kernel 调用的细节将在后面的章节中解释。尖括号内的值表示我们希望在运行时从宿主传递到设备的参数。基本上，它表示将在设备上并行运行的块和线程的数量。因此，在这个代码中，<< <1,1> >> 表示 myfirstkernel 将在设备上的一个块和一个线程或块上运行。尽管这不是设备资源的最佳使用，但它是一个理解在宿主上执行和在设备上执行的代码之间差异的好起点。

再次，为了回顾和修改 Hello, CUDA! 代码，myfirstkernel 函数将在一个块和一个线程或块上运行在设备上。它将通过在主函数内部的宿主代码中调用一种称为 kernel launch 的方法来启动。

编写代码后，你将如何执行此代码并查看输出？下一节将描述在 Windows 和 Ubuntu 上编写和执行 Hello, CUDA! 代码的步骤。

在 Windows 上创建 CUDA C 程序的步骤

本节描述了在 Windows 上使用 Visual Studio 创建和执行基本 CUDA C 程序的步骤。步骤如下：

打开 Microsoft Visual Studio。
前往文件 | 新建 | 工程。
选择 NVIDIA | CUDA 9.0 | CUDA 9.0 Runtime。
给项目命名你想要的名称，然后点击确定。
它将创建一个包含示例 kernel.cu 文件的工程。现在通过双击它来打开此文件。
从文件中删除现有的代码，并写入之前给出的代码。
从“构建”选项卡构建项目，并按 Ctrl + F5 调试代码。如果一切正常，你将在命令行看到 Hello, CUDA! 如此显示：

在 Ubuntu 上创建 CUDA C 程序的步骤

本节描述了在 Ubuntu 上使用 Nsight Eclipse 插件创建和执行基本 CUDA C 程序的步骤。步骤如下：

通过在终端中打开终端并输入 nsight 来打开 Nsight。
前往文件 | 新建 | CUDA C/C++ 项目。
给项目起一个你喜欢的名字，然后点击确定。
它将创建一个带有示例文件的项目。现在通过双击它来打开此文件。
从文件中删除现有的代码，并写入之前给出的代码。
通过按播放按钮运行代码。如果一切正常，你将在终端看到 Hello, CUDA! 如此显示：

摘要

总结来说，在本章中，你被介绍了 CUDA 并简要介绍了并行计算的重要性。详细讨论了 CUDA 和 GPU 在各个领域的应用。本章描述了在个人电脑上执行 CUDA 应用所需的硬件和软件设置。它提供了一个逐步的过程，用于在本地电脑上安装 CUDA。

最后一节通过开发一个简单的程序并在 Windows 和 Ubuntu 上执行它，为 CUDA C 应用程序开发提供了一个入门指南。

在下一章中，我们将基于 CUDA C 编程的知识来构建。你将通过几个实际示例了解如何使用 CUDA C 进行并行计算，以展示它相对于常规编程的快速性。你还将了解线程和块的概念以及如何在多个线程和块之间进行同步。

问题

解释三种提高计算硬件性能的方法。哪种方法用于开发 GPU？
正误判断：提高延迟将提高吞吐量。
填空：CPU 设计用于提高 ___，而 GPU 设计用于提高 ___。
以从一个地方到另一个地方旅行为例，距离为 240 公里。你可以选择一辆可以容纳五人的汽车，速度为 60 公里/小时，或者一辆可以容纳 40 人的公共汽车，速度为 40 公里/小时。哪种选项将提供更好的延迟，哪种选项将提供更好的吞吐量？
解释使 GPU 和 CUDA 在计算机视觉应用中特别有用的原因。
正误判断：CUDA 编译器不能编译没有设备代码的代码。
在本章讨论的 Hello, CUDA! 示例中，printf 语句将由主机还是设备执行？

第二章：使用 CUDA C 进行并行编程

在上一章中，我们看到了安装 CUDA 和使用它的程序是多么容易。尽管示例并不令人印象深刻，但它被用来向您证明开始使用 CUDA 非常容易。在本章中，我们将在此基础上进行构建。它详细介绍了如何使用 CUDA 为 GPU 编写高级程序。它从一个变量加法程序开始，然后逐步构建到 CUDA C 中的复杂向量操作示例。它还涵盖了内核的工作原理以及如何在 CUDA 程序中使用设备属性。本章讨论了在 CUDA 程序中如何操作向量，以及与 CUDA 编程相关的术语。

本章将涵盖以下主题：

内核调用的概念
在 CUDA 中创建内核函数并将参数传递给它
配置 CUDA 程序的内核参数和内存分配
CUDA 程序中的线程执行
从 CUDA 程序中访问 GPU 设备属性
在 CUDA 程序中处理向量
并行通信模式

技术要求

本章需要熟悉基本的 C 或 C++ 编程语言，特别是动态内存分配函数。本章中使用的所有代码都可以从以下 GitHub 链接下载：github.com/PacktPublishing/Hands-On-GPU-Accelerated-Computer-Vision-with-OpenCV-and-CUDA。代码可以在任何操作系统上执行，尽管它只在 Windows 10 和 Ubuntu 16.04 上进行了测试。

查看以下视频以查看代码的实际运行情况：

bit.ly/2PQmu4O

CUDA 程序结构

我们之前已经看到了一个非常简单的 Hello, CUDA! 程序，它展示了与 CUDA 程序相关的一些重要概念。CUDA 程序是由在主机或 GPU 设备上执行的功能组合而成。不显示并行性的功能在 CPU 上执行，而显示数据并行的功能在 GPU 上执行。GPU 编译器在编译过程中将这些功能分开。正如前一章所看到的，用于在设备上执行的功能使用 __global__ 关键字定义，并由 NVCC 编译器编译，而正常的 C 主机代码由 C 编译器编译。CUDA 代码基本上是相同的 ANSI C 代码，增加了用于利用数据并行的某些关键字。

因此，在本节中，我们通过一个简单的双变量加法程序来解释与 CUDA 编程相关的重要概念，例如内核调用、从主机到设备传递内核函数的参数、内核参数的配置、用于利用数据并行的 CUDA API，以及主机和设备上的内存分配是如何进行的。

CUDA C 中的双变量加法程序

在第一章中看到的简单 Hello, CUDA! 代码，介绍 CUDA 和开始使用 CUDA，设备函数是空的。它没有任何作用。本节解释了一个简单的加法程序，该程序在设备上执行两个变量的加法。尽管它没有利用设备的任何数据并行性，但它对于展示 CUDA C 的重要编程概念非常有用。首先，我们将看到如何编写一个用于加法两个变量的内核函数。

内核函数的代码如下所示：

include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>
//Definition of kernel function to add two variables
__global__ void gpuAdd(int d_a, int d_b, int *d_c) 
{
   *d_c = d_a + d_b;
}

gpuAdd 函数看起来与在 ANSI C 中实现的正常 add 函数非常相似。它接受两个整数变量 d_a 和 d_b 作为输入，并将加法存储在由第三个整数指针 d_c 指示的内存位置。设备函数的返回值是 void，因为它将答案存储在设备指针指向的内存位置，而不是显式返回任何值。现在我们将看到如何编写此代码的主函数。主函数的代码如下所示：


 int main(void) 
{
 //Defining host variable to store answer
   int h_c;
 //Defining device pointer
   int *d_c;
 //Allocating memory for device pointer
   cudaMalloc((void**)&d_c, sizeof(int));
 //Kernel call by passing 1 and 4 as inputs and storing answer in d_c
 //<< <1,1> >> means 1 block is executed with 1 thread per block
   gpuAdd << <1, 1 >> > (1, 4, d_c);
 //Copy result from device memory to host memory
   cudaMemcpy(&h_c, d_c, sizeof(int), cudaMemcpyDeviceToHost);
   printf("1 + 4 = %d\n", h_c);
 //Free up memory
   cudaFree(d_c);
   return 0;
}

在 main 函数中，前两行定义了主机和设备的变量。第三行使用 cudaMalloc 函数在设备上为 d_c 变量分配内存。cudaMalloc 函数与 C 中的 malloc 函数类似。在主函数的第四行中，gpuAdd 被调用，使用 1 和 4 作为两个输入变量，以及 d_c，这是一个设备内存指针，作为输出指针变量。gpuAdd 函数的奇怪语法，也称为内核调用，将在下一节中解释。如果需要将 gpuAdd 的答案用于主机，则必须通过 cudaMemcpy 函数将其从设备的内存复制到主机的内存。然后，使用 printf 函数打印此答案。最后一行使用 cudafree 函数释放设备上使用的内存。显式释放程序中使用的所有设备内存非常重要；否则，你可能会在某个时刻耗尽内存。以 // 开头的行是用于提高代码可读性的注释，这些行被编译器忽略。

双变量加法程序有两个函数，main 和 gpuAdd。正如你所见，gpuAdd 是通过使用 __global__ 关键字定义的，因此它旨在在设备上执行，而主函数将在主机上执行。该程序在设备上添加两个变量，并在命令行上打印输出，如下所示：

在本书中，我们将使用一个约定，即主机变量将以前缀 h_ 开头，设备变量将以前缀 d_ 开头。这不是强制性的；这只是为了让读者能够轻松理解概念，而不会在主机和设备之间产生混淆。

所有 CUDA API，如 cudaMalloc、cudaMemcpy 和 cudaFree，以及其他重要的 CUDA 编程概念，如内核调用、向内核传递参数和内存分配问题，将在接下来的章节中讨论。

内核调用

使用 ANSI C 关键字以及 CUDA 扩展关键字编写的设备代码被称为内核。它通过一种称为 内核调用 的方法从主机代码中启动。基本上，内核调用的意义是我们从主机代码中启动设备代码。内核调用通常会产生大量的块和线程，以在 GPU 上利用数据并行性。内核代码与普通 C 函数非常相似；只是这些代码是由多个线程并行执行的。它有一个非常奇怪的语法，如下所示：

kernel << <number of blocks, number of threads per block, size of shared memory > >> (parameters for kernel)

它以我们想要启动的内核的名称开始。你应该确保这个内核是用 __global__ 关键字定义的。然后，它有 << < > >> 内核启动运算符，其中包含内核的配置参数。它可以包含三个用逗号分隔的参数。第一个参数表示你想要执行的块的数量，第二个参数表示每个块将有多少线程。因此，内核启动启动的总线程数将是这两个数字的乘积。第三个参数，指定内核使用的共享内存的大小，是可选的。在变量加法程序中，内核启动的语法如下：

gpuAdd << <1,1> >> (1 , 4, d_c)

在这里，gpuAdd 是我们想要启动的内核的名称，而 <<<1,1>>> 表示我们想要每个块有一个线程，这意味着我们只启动了一个线程。圆括号中的三个参数是传递给内核的参数。在这里，我们传递了两个常量，1 和 4。第三个参数是指向设备内存 d_c 的指针。它指向内核在加法操作后将在设备内存中存储答案的位置。程序员必须注意的一点是，传递给内核的指针应该只指向设备内存。如果它指向主机内存，可能会使你的程序崩溃。内核执行完成后，设备指针指向的结果可以被复制回主机内存以供进一步使用。在设备上仅启动一个线程进行执行并不是设备资源的最佳使用方式。假设你想要并行启动多个线程；你需要在内核调用的语法中进行哪些修改？这将在下一节中讨论，并被称为“配置内核参数”。

配置内核参数

为了在设备上并行启动多个线程，我们不得不在内核调用中配置参数，这些参数被写入内核启动操作符内部。它们指定了每个块中的线程数和块的数量。我们可以通过每个块中的多个线程并行启动多个块。通常，每个块中的线程数限制为 512 或 1,024。每个块在流式多处理器上运行，一个块中的线程可以通过共享内存相互通信。程序员无法选择哪个多处理器将执行特定的块，以及块或线程执行的顺序。

假设你想并行启动 500 个线程；你可以对之前显示的内核启动语法进行哪些修改？一个选项是通过以下语法启动一个包含 500 个线程的块：

gpuAdd<< <1,500> >> (1,4, d_c)

我们也可以启动 500 个每个线程的块或 250 个线程的 2 个块。相应地，你必须修改内核启动操作符中的值。程序员必须小心，确保每个块中的线程数不超过 GPU 设备的最大支持限制。在这本书中，我们针对计算机视觉应用，需要处理二维和三维图像。在这里，如果块和线程不是一维的，而是更多维的，将有助于更好的处理和可视化。

GPU 支持三维块网格和三维线程块。它具有以下语法：

mykernel<< <dim3(Nbx, Nby,Nbz), dim3(Ntx, Nty,Ntz) > >> ()

这里 N[bx]、N[by] 和 N[bz] 分别表示在 x、y 和 z 轴方向上的网格中的块数。同样，N[t][x]、N[ty] 和 N[tz] 表示在 x、y 和 z 轴方向上的块中的线程数。如果 y 和 z 维度未指定，则默认为 1。例如，为了处理图像，你可以启动一个 16 x 16 的块网格，所有块都包含 16 x 16 的线程。语法如下：

mykernel << <dim3(16,16),dim3(16,16)> >> ()

总结来说，在启动内核时配置块和线程的数量非常重要。应根据我们正在处理的应用程序和 GPU 资源进行适当的考虑。下一节将解释一些在常规 ANSI C 函数之上添加的重要 CUDA 函数。

CUDA API 函数

在变量加法程序中，我们遇到了一些对常规 C 或 C++程序员来说不熟悉的函数或关键字。这些关键字和函数包括 __global__、cudaMalloc、cudaMemcpy 和 cudaFree。因此，在本节中，我们将逐一详细解释这些函数：

global：这是三个限定符关键字之一，与 __device__ 和 __host__ 一起。此关键字表示一个函数被声明为设备函数，当从主机调用时将在设备上执行。请注意，此函数只能从主机调用。如果您想使您的函数在设备上执行并从设备函数调用，那么您必须使用 __device__ 关键字。__host__ 关键字用于定义只能从其他主机函数调用的主机函数。这类似于正常的 C 函数。默认情况下，程序中的所有函数都是主机函数。__host__ 和 __device__ 可以同时使用来定义任何函数。它将生成相同函数的两个副本。一个将在主机上执行，另一个将在设备上执行。
cudaMalloc：它类似于 C 中用于动态内存分配的 Malloc 函数。此函数用于在设备上分配特定大小的内存块。以下是一个 cudaMalloc 的语法示例：

cudaMalloc(void ** d_pointer, size_t size)
Example: cudaMalloc((void**)&d_c, sizeof(int));

如前一个示例代码所示，它分配了一个与一个整型变量大小相等的内存块，并返回指向该内存位置的指针 d_c。

cudaMemcpy：此函数类似于 C 中的 Memcpy 函数。它用于将一块内存从一个主机或设备上的其他块复制。它具有以下语法：

cudaMemcpy ( void * dst_ptr, const void * src_ptr, size_t size, enum cudaMemcpyKind kind )
Example: cudaMemcpy(&h_c, d_c, sizeof(int), cudaMemcpyDeviceToHost);

此函数有四个参数。前两个参数是目标指针和源指针，它们指向主机或设备内存位置。第三个参数指示复制的尺寸，最后一个参数指示复制的方向。可以是主机到设备、设备到设备、主机到主机或设备到主机。但请注意，您必须将此方向与前两个参数中的适当指针匹配。如示例所示，我们通过指定设备指针 d_c 为源，主机指针 h_c 为目标，将一个整型变量的块从设备复制到主机。

cudaFree：它类似于 C 中可用的 free 函数。cudaFree 的语法如下：

cudaFree ( void * d_ptr )
Example: cudaFree(d_c)

它释放由 d_ptr 指向的内存空间。在示例代码中，它释放了由 d_c 指向的内存位置。请确保使用 cudaMalloc 分配了 d_c 的内存，然后使用 cudaFree 释放它。

在 CUDA 中，除了现有的 ANSI C 函数之外，还有很多其他关键字和函数可用。我们将经常使用这三个函数，因此它们在本节中进行了讨论。更多详情，您可以随时访问 CUDA 编程指南。

向 CUDA 函数传递参数

变量加法程序的 gpuAdd 内核函数与正常的 C 函数非常相似。因此，像正常的 C 函数一样，内核函数也可以按值或按引用传递参数。因此，在本节中，我们将看到传递 CUDA 内核参数的两种方法。

按值传递参数

如果你还记得，在 gpuAdd 程序中，调用内核的语法如下：

gpuAdd << <1,1> >>(1,4,d_c)

另一方面，gpuAdd 函数在定义中的签名如下：

__global__  gpuAdd(int d_a, int d_b, int *d_c)

因此，你可以看到我们在调用内核时传递了 d_a 和 d_b 的值。首先，参数 1 将被复制到 d_a，然后参数 4 将在调用内核时复制到 d_b。加法操作后的答案将存储在设备内存中由 d_c 指向的地址。我们也可以直接将值 1 和 4 作为内核的输入，如下所示：

gpuAdd << <1,1> >>(a,b,d_c)

在这里，a 和 b 是可以包含任何整数值的整型变量。按值传递参数不建议使用，因为它会在程序中造成不必要的混淆和复杂化。最好是通过引用传递参数。

按引用传递参数

现在我们将看到如何通过引用传递参数来编写相同的程序。为此，我们首先需要修改用于两个变量加法的内核函数。按引用传递参数的修改后的内核如下所示：

#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>
//Kernel function to add two variables, parameters are passed by reference
 __global__ void gpuAdd(int *d_a, int *d_b, int *d_c) 
{
  *d_c = *d_a + *d_b;
}

在将整数变量 d_a 和 d_b 作为内核的输入时，我们取这些变量在设备上的指针 *d_a 和 *d_b 作为输入。加法操作后的答案将存储在由第三个整数指针 d_c 指向的内存位置。传递给这个设备函数的指针应该使用 cudaMalloc 函数分配内存。此代码的主函数如下所示：

int main(void) 
{
  //Defining host and variables
  int h_a,h_b, h_c;
  int *d_a,*d_b,*d_c;
  //Initializing host variables
  h_a = 1;
  h_b = 4;
  //Allocating memory for Device Pointers
  cudaMalloc((void**)&d_a, sizeof(int));
  cudaMalloc((void**)&d_b, sizeof(int));
  cudaMalloc((void**)&d_c, sizeof(int));
  //Coping value of host variables in device memory
  cudaMemcpy(d_a, &h_a, sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_b, &h_b, sizeof(int), cudaMemcpyHostToDevice);
  //Calling kernel with one thread and one block with parameters passed by reference
  gpuAdd << <1, 1 >> > (d_a, d_b, d_c);
  //Coping result from device memory to host
  cudaMemcpy(&h_c, d_c, sizeof(int), cudaMemcpyDeviceToHost);
  printf("Passing Parameter by Reference Output: %d + %d = %d\n", h_a, h_b, h_c);
  //Free up memory
  cudaFree(d_a);
  cudaFree(d_b);
  cudaFree(d_c);
  return 0;
 }

h_a、h_b 和 h_c 是主机内存中的变量。它们像正常的 C 代码一样定义。另一方面，d_a、d_b 和 d_c 是位于主机内存中的指针，它们指向设备内存。它们通过使用 cudaMalloc 函数从主机分配内存。h_a 和 h_b 的值通过使用 cudaMemcpy 函数复制到由 d_a 和 d_b 指向的设备内存中，数据传输方向是从主机到设备。然后，在内核调用中，这三个设备指针作为参数传递给内核。内核执行加法操作并将结果存储在由 d_c 指向的内存位置。结果再次通过 cudaMemcpy 复制回主机内存，但这次数据传输方向是从设备到主机。程序输出如下：

程序结束时使用 cudaFree 释放三个设备指针使用的内存。主机和设备上的示例内存映射将类似于以下内容：

主机内存（CPU）	设备内存（GPU）
地址	值
#01	h_a=1
#02	h_b=4
#03	h_c=5
#04	d_a=#01
#05	d_b=#02
#06	d_c=#03

从表中可以看出，d_a、d_b和d_c位于主机上，并指向设备内存中的值。在通过引用传递参数给内核时，你应该注意所有指针都只指向设备内存。如果不是这样，程序可能会崩溃。

在使用设备指针并将它们传递给内核时，程序员必须遵循一些限制。使用cudaMalloc分配内存的设备指针只能用于从设备内存中读取或写入。它们可以作为参数传递给设备函数，但不应用于从主机函数中读取和写入内存。为了简化，设备指针应用于从设备函数中读取和写入设备内存，而主机指针应用于从主机函数中读取和写入主机内存。因此，在本书中，你将始终在内核函数中看到以d_为前缀的设备指针。

总结来说，在本节中，通过以两个变量的附加程序为例，详细解释了与 CUDA 编程相关的概念。在本节之后，你应该熟悉基本的 CUDA 编程概念以及与 CUDA 程序相关的术语。在下一节中，你将学习如何在设备上执行线程。

在设备上执行线程

我们已经看到，在配置内核参数时，我们可以并行启动多个块和多个线程。那么，这些块和线程的启动和完成执行的顺序是怎样的呢？如果我们想在其他线程中使用一个线程的输出，了解这一点很重要。为了理解这一点，我们修改了第一章节中看到的hello,CUDA!程序中的内核，通过在内核调用中包含一个打印语句来打印块号。修改后的代码如下：

#include <iostream>
#include <stdio.h>
__global__ void myfirstkernel(void) 
{
  //blockIdx.x gives the block number of current kernel
   printf("Hello!!!I'm thread in block: %d\n", blockIdx.x);
}
int main(void) 
{
   //A kernel call with 16 blocks and 1 thread per block
   myfirstkernel << <16,1>> >();

   //Function used for waiting for all kernels to finish
   cudaDeviceSynchronize();

   printf("All threads are finished!\n");
   return 0;
}

从代码中可以看出，我们正在并行启动一个内核，有 16 个块，每个块有一个线程。在内核代码中，我们正在打印内核执行的块 ID。我们可以认为 16 个相同的myfirstkernel副本并行开始执行。每个副本都将有一个唯一的块 ID，可以通过blockIdx.x CUDA指令访问，以及一个唯一的线程 ID，可以通过threadIdx.x访问。这些 ID 将告诉我们哪个块和线程正在执行内核。当你多次运行程序时，你会发现，每次块执行的顺序都不同。一个示例输出如下所示：

您应该问的一个问题是，前面的程序将产生多少种不同的输出模式？正确的答案是 16！它将产生n阶乘数量的输出，其中n表示并行启动的块的数量。因此，每次在 CUDA 中编写程序时，您都应该小心，确保块以随机顺序执行。

此程序还包含一个额外的 CUDA 指令：cudaDeviceSynchronize()。为什么使用它？这是因为内核启动是一个异步过程，这意味着在内核完成执行之前，它会立即将控制权返回给启动 GPU 进程之前的 CPU 线程。在前面的代码中，CPU 线程的下一行是print，应用程序退出将在内核完成执行之前终止控制台。因此，如果我们不包括此指令，您将看不到任何内核执行的打印语句。内核随后生成的输出将无处可去，您将看不到它。为了看到内核生成的输出，我们将包括此指令，这确保了内核在应用程序被允许退出之前完成，并且内核的输出将找到等待的标准输出队列。

从 CUDA 程序中访问 GPU 设备属性

CUDA 提供了一个简单的接口来查找信息，例如确定哪些 CUDA 启用型 GPU 设备（如果有）存在以及每个设备支持哪些功能。首先，重要的是要获取系统上 CUDA 启用型设备数量的统计，因为一个系统可能包含多个启用 GPU 的设备。这个数量可以通过 CUDA API cudaGetDeviceCount() 来确定。获取系统上 CUDA 启用型设备数量的程序如下所示：

#include <memory>
#include <iostream>
#include <cuda_runtime.h>
// Main Program 
int main(void)
{
  int device_Count = 0;
  cudaGetDeviceCount(&device_Count);
  // This function returns count of number of CUDA enable devices and 0 if there are no CUDA capable devices.
  if (device_Count == 0)
  {
     printf("There are no available device(s) that support CUDA\n");
  }
  else
  {
     printf("Detected %d CUDA Capable device(s)\n", device_Count);
  }
}

通过查询cudaDeviceProp结构可以找到每个设备的相关信息，该结构返回所有设备属性。如果您有多个 CUDA 能力型设备，则可以启动一个 for 循环来遍历所有设备属性。以下部分包含设备属性列表，分为不同的集合以及用于从 CUDA 程序中访问它们的简短代码片段。这些属性由 CUDA 9 运行时中的cudaDeviceProp结构提供。

如需了解 CUDA 不同版本中属性的相关详细信息，您可以查看特定版本的编程指南。

通用设备属性

cudaDeviceProp提供了几个属性，可用于识别设备和正在使用的版本。它提供了一个name属性，该属性以字符串形式返回设备名称。我们还可以通过查询cudaDriverGetVersion和cudaRuntimeGetVersion属性来获取设备使用的驱动程序和运行时引擎的版本。有时，如果您有多个设备，您可能希望使用具有更多多处理器的设备。multiProcessorCount属性返回设备上多处理器的数量。通过使用clockRate属性可以获取 GPU 的时钟速度，它以千赫兹为单位返回时钟频率。以下代码片段展示了如何从 CUDA 程序中使用这些属性：

cudaDeviceProp device_Property;
cudaGetDeviceProperties(&device_Property, device);
printf("\nDevice %d: \"%s\"\n", device, device_Property.name);
cudaDriverGetVersion(&driver_Version);
cudaRuntimeGetVersion(&runtime_Version);
printf(" CUDA Driver Version / Runtime Version %d.%d / %d.%d\n", driver_Version / 1000, (driver_Version % 100) / 10, runtime_Version / 1000, (runtime_Version % 100) / 10);
printf( " Total amount of global memory: %.0f MBytes (%llu bytes)\n",
 (float)device_Property.totalGlobalMem / 1048576.0f, (unsigned long long) device_Property.totalGlobalMem);
 printf(" (%2d) Multiprocessors", device_Property.multiProcessorCount );
printf("  GPU Max Clock rate: %.0f MHz (%0.2f GHz)\n", device_Property.clockRate * 1e-3f, device_Property.clockRate * 1e-6f);

与内存相关的属性

GPU 上的内存具有分层架构。它可以按 L1 缓存、L2 缓存、全局内存、纹理内存和共享内存来划分。cudaDeviceProp提供了许多属性，有助于识别设备上可用的内存。memoryClockRate和memoryBusWidth分别提供内存的时钟频率和总线宽度。内存的速度非常重要，它会影响您程序的整体速度。totalGlobalMem返回设备上可用的全局内存大小。totalConstMem返回设备上可用的总常量内存。sharedMemPerBlock返回设备中可以使用的总共享内存。每个块可用的寄存器总数可以通过使用regsPerBlock来识别。L2 缓存的大小可以通过l2CacheSize属性来识别。以下代码片段展示了如何从 CUDA 程序中使用内存相关的属性：

printf( " Total amount of global memory: %.0f MBytes (%llu bytes)\n",
(float)device_Property.totalGlobalMem / 1048576.0f, (unsigned long long) device_Property.totalGlobalMem);
printf(" Memory Clock rate: %.0f Mhz\n", device_Property.memoryClockRate * 1e-3f);
printf(" Memory Bus Width: %d-bit\n", device_Property.memoryBusWidth);
if (device_Property.l2CacheSize)
{
    printf(" L2 Cache Size: %d bytes\n", device_Property.l2CacheSize);
}
printf(" Total amount of constant memory: %lu bytes\n",         device_Property.totalConstMem);
printf(" Total amount of shared memory per block: %lu bytes\n", device_Property.sharedMemPerBlock);
printf(" Total number of registers available per block: %d\n", device_Property.regsPerBlock);

与线程相关的属性

如前几节所示，块和线程可以是多维的。因此，了解每个维度中可以并行启动多少线程和块将很有帮助。每个多处理器和每个块的线程数量也有上限。这个数量可以通过使用maxThreadsPerMultiProcessor和maxThreadsPerBlock来找到。这在内核参数配置中非常重要。如果您在每个块中启动的线程数超过了每个块可能的最大线程数，则您的程序可能会崩溃。每个维度中每个块的最大线程数可以通过maxThreadsDim来识别。同样，每个维度中每个网格的最大块数可以通过使用maxGridSize来识别。这两个属性都返回一个包含三个值的数组，分别表示x、y和z维度上的最大值。以下代码片段展示了如何从 CUDA 代码中使用线程相关的属性：

printf(" Maximum number of threads per multiprocessor: %d\n",              device_Property.maxThreadsPerMultiProcessor);
printf(" Maximum number of threads per block: %d\n",         device_Property.maxThreadsPerBlock);
printf(" Max dimension size of a thread block (x,y,z): (%d, %d, %d)\n",
    device_Property.maxThreadsDim[0],
    device_Property.maxThreadsDim[1],
    device_Property.maxThreadsDim[2]);
printf(" Max dimension size of a grid size (x,y,z): (%d, %d, %d)\n",
    device_Property.maxGridSize[0],
    device_Property.maxGridSize[1],
    device_Property.maxGridSize[2]);

在cudaDeviceProp结构体中还有许多其他属性可用。您可以查阅 CUDA 编程指南以获取其他属性的详细信息。以下是在 NVIDIA Geforce 940MX GPU 和 CUDA 9.0 上执行并组合所有先前代码段输出的结果：

你可能会问的一个问题是，为什么你应该对了解设备属性感兴趣。答案是，这将帮助你在存在多个 GPU 设备的情况下选择具有更多多处理器的 GPU 设备。如果您的应用程序中的内核需要与 CPU 进行紧密交互，那么您可能希望内核在共享系统内存的集成 GPU 上运行。这些属性还将帮助您找到设备上可用的块数和每个块中的线程数。这将帮助您配置内核参数。为了向您展示设备属性的一个用途，假设您有一个需要双精度浮点运算的应用程序。并非所有 GPU 设备都支持此操作。为了知道您的设备是否支持双精度浮点运算，并将该设备设置为您的应用程序，可以使用以下代码：

#include <memory>
#include <iostream>
#include <cuda_runtime.h>
// Main Program
int main(void)
{
int device;
cudaDeviceProp device_property;
cudaGetDevice(&device);
printf("ID of device: %d\n", device);
memset(&device_property, 0, sizeof(cudaDeviceProp));
device_property.major = 1;
device_property.minor = 3;
cudaChooseDevice(&device, &device_property);
printf("ID of device which supports double precision is: %d\n", device);
cudaSetDevice(device);
}

此代码使用了cudaDeviceprop结构中可用的两个属性，这些属性有助于确定设备是否支持双精度运算。这两个属性是主版本号和次版本号。CUDA 文档告诉我们，如果主版本号大于 1 且次版本号大于 3，则该设备将支持双精度运算。因此，程序中的device_property结构被填充了这两个值。CUDA 还提供了cudaChooseDevice API，该 API 有助于选择具有特定属性的设备。此 API 用于当前设备，以确定它是否包含这两个属性。如果包含属性，则使用cudaSetDevice API 选择该设备用于您的应用程序。如果系统中存在多个设备，则此代码应编写在一个循环中，以便遍历所有设备。

虽然很简单，但这一部分对于您了解哪些应用程序可以由您的 GPU 设备支持以及哪些不支持非常重要。

CUDA 中的向量运算

到目前为止，我们看到的程序都没有利用 GPU 设备的并行处理能力。它们只是编写来让您熟悉 CUDA 中的编程概念。从本节开始，我们将通过在 GPU 上执行向量或数组运算来利用 GPU 的并行处理能力。

两个向量加法程序

要理解 GPU 上的向量运算，我们首先将在 CPU 上编写一个向量加法程序，然后修改它以利用 GPU 的并行结构。我们将取两个数字数组，并将逐元素加法的结果存储在第三个数组中。CPU 上的向量加法函数如下所示：

#include "stdio.h"
#include<iostream>
 //Defining Number of elements in Array
#define N 5
 //Defining vector addition function for CPU
void cpuAdd(int *h_a, int *h_b, int *h_c) 
{
     int tid = 0;
     while (tid < N)
     {
         h_c[tid] = h_a[tid] + h_b[tid];
         tid += 1;
     }
 }

cpuAdd 应该非常容易理解。你可能觉得难以理解的是 tid 的使用。它被包含进来是为了使程序与 GPU 程序相似，其中 tid 表示特定的线程 ID。在这里，如果你有一个多核 CPU，那么你可以为每个核心初始化 tid 为 0 和 1，然后在循环中将其加 2，这样其中一个 CPU 将对偶数元素进行求和，而另一个 CPU 将对奇数元素进行加法。代码的 main 函数如下所示：

int main(void) 
{
   int h_a[N], h_b[N], h_c[N];
   //Initializing two arrays for addition
   for (int i = 0; i < N; i++) 
   {
     h_a[i] = 2 * i*i;
     h_b[i] = i;
     }
   //Calling CPU function for vector addition
   cpuAdd (h_a, h_b, h_c);
   //Printing Answer
   printf("Vector addition on CPU\n");
   for (int i = 0; i < N; i++) 
   {
     printf("The sum of %d element is %d + %d = %d\n", i, h_a[i], h_b[i],             h_c[i]);
   }
   return 0;
 }

程序中有两个函数：main 和 cpuAdd。在 main 函数中，我们首先定义了两个数组来存储输入，并将其初始化为一些随机数。然后，我们将这两个数组作为输入传递给 cpuAdd 函数。cpuAdd 函数将答案存储在第三个数组中。然后，我们在控制台上打印这个答案，如下所示：

使用 tid in cpuadd 函数的解释可能给你一些如何为 GPU 执行编写相同函数的思路，因为 GPU 可以并行处理多个核心。如果我们用那个核心的 ID 来初始化这个加法函数，那么我们就可以并行地对所有元素进行加法运算。因此，GPU 上加法操作的修改后的内核函数如下所示：

#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
 //Defining number of elements in Array
#define N 5
 //Defining Kernel function for vector addition
__global__ void gpuAdd(int *d_a, int *d_b, int *d_c) 
{
 //Getting block index of current kernel
     int tid = blockIdx.x; // handle the data at this index
     if (tid < N)
     d_c[tid] = d_a[tid] + d_b[tid];
 }

在 gpuAdd 内核函数中，tid 被初始化为当前内核执行的块的块 ID。所有内核都将添加由该块 ID 索引的数组元素。如果块的数量等于数组中的元素数量，那么所有加法操作都将并行执行。接下来将解释如何从 main 函数中调用这个内核。main 函数的代码如下：

int main(void) 
{
 //Defining host arrays
 int h_a[N], h_b[N], h_c[N];
 //Defining device pointers
 int *d_a, *d_b, *d_c;
 // allocate the memory
 cudaMalloc((void**)&d_a, N * sizeof(int));
 cudaMalloc((void**)&d_b, N * sizeof(int));
 cudaMalloc((void**)&d_c, N * sizeof(int));
 //Initializing Arrays
 for (int i = 0; i < N; i++) 
    {
     h_a[i] = 2*i*i;
     h_b[i] = i ;
     }

// Copy input arrays from host to device memory
 cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice);
 cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice);

//Calling kernels with N blocks and one thread per block, passing device pointers as parameters
gpuAdd << <N, 1 >> >(d_a, d_b, d_c);
 //Copy result back to host memory from device memory
cudaMemcpy(h_c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);
printf("Vector addition on GPU \n");
 //Printing result on console
for (int i = 0; i < N; i++) 
{
     printf("The sum of %d element is %d + %d = %d\n", i, h_a[i], h_b[i],             h_c[i]);
}
 //Free up memory
 cudaFree(d_a);
 cudaFree(d_b);
 cudaFree(d_c);
 return 0;
}

GPU 的 main 函数具有本章第一部分所述的已知结构：

它从定义主机和设备的数组和指针开始。使用 cudaMalloc 函数为设备指针分配内存。
要传递给内核的数组通过使用 cudaMemcpy 函数从主机内存复制到设备内存。
内核是通过将设备指针作为参数传递给它来启动的。如果你看到内核启动操作符内的值，它们是 N 和 1，这表示我们正在启动 N 个块，每个块有一个线程。
内核在设备内存中存储的答案通过再次使用 cudaMemcpy 被复制回主机内存，但这次数据传输的方向是从设备到主机。
最后，使用 cudaFree 函数释放分配给三个设备指针的内存。程序的输出如下：

所有 CUDA 程序都遵循之前显示的相同模式。我们在并行启动 N 个块。这意味着我们同时启动了 N 个相同的内核副本。你可以通过一个现实生活中的例子来理解这一点：假设你想要将五个大箱子从一个地方运到另一个地方。在第一种方法中，你可以通过雇佣一个人来完成这个任务，这个人从一处运到另一处，然后重复五次。这个选项会花费时间，这类似于向量在 CPU 上是如何加的。现在，假设你雇佣了五个人，每个人携带一个箱子。他们每个人也知道他们携带的箱子的 ID。这个选项将比之前的选项快得多。他们每个人只需要被告知他们必须携带一个特定 ID 的箱子从一处运到另一处。

这正是内核在设备上定义和执行的方式。每个内核副本都知道自己的 ID。这可以通过blockIdx.x命令来知道。每个副本在其 ID 索引的数组元素上工作。所有副本并行地添加所有元素，这显著减少了整个数组的处理时间。所以，从某种意义上说，我们通过在 CPU 的顺序执行上并行执行操作来提高吞吐量。CPU 代码和 GPU 代码之间的吞吐量比较将在下一节中解释。

比较 CPU 和 GPU 代码之间的延迟

CPU 和 GPU 的加法程序以模块化的方式编写，这样你可以玩转 N 的值。如果 N 很小，那么你不会注意到 CPU 和 GPU 代码之间有显著的时间差异。但是，如果你 N 足够大，那么你将注意到相同向量加法中 CPU 执行时间和 GPU 执行时间的显著差异。可以通过在现有代码中添加以下行来测量特定块的执行时间：

clock_t start_d = clock();
printf("Doing GPU Vector add\n");
gpuAdd << <N, 1 >> >(d_a, d_b, d_c);
cudaThreadSynchronize();
clock_t end_d = clock();
double time_d = (double)(end_d - start_d) / CLOCKS_PER_SEC;
printf("No of Elements in Array:%d \n Device time %f seconds \n host time %f Seconds\n", N, time_d, time_h);

时间是通过计算执行特定操作所花费的总时钟周期数来衡量的。这可以通过使用clock()函数测量的开始和结束的时钟滴答计数之差来完成。这个差值除以每秒的时钟周期数，以得到执行时间。当在 CPU 和 GPU 之前的向量加法程序中将 N 设置为 10,000,000 并同时执行时，输出如下：

从输出中可以看出，当相同的函数在 GPU 上实现时，执行时间或吞吐量从 25 毫秒提高到了几乎 1 毫秒。这证明了我们在理论中之前看到的事实，即在 GPU 上并行执行代码有助于提高吞吐量。CUDA 提供了一个高效且准确的方法来测量 CUDA 程序的性能，使用 CUDA 事件，这将在后面的章节中解释。

CUDA 中向量的逐元素平方

现在，你可以问的一个问题是，既然我们正在每个块中用一个线程并行启动 N 个块，我们能否以相反的方式工作？答案是是的。我们可以并行地只启动一个包含 N 个线程的块。为了展示这一点并让你更熟悉在 CUDA 中围绕向量工作，我们以数组中数字逐元素平方的第二个例子为例。我们取一个数字数组，并返回一个包含这些数字平方的数组。用于找到逐元素平方的内核函数如下所示：

#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
 //Defining number of elements in Array
#define N 5
//Kernel function for squaring number
__global__ void gpuSquare(float *d_in, float *d_out) 
{
     //Getting thread index for current kernel
     int tid = threadIdx.x; // handle the data at this index
     float temp = d_in[tid];
     d_out[tid] = temp*temp;
 }

gpuSquare 内核函数有两个数组的指针作为参数。第一个指针 d_in 指向存储输入数组的内存位置，而第二个指针 d_out 指向存储输出的内存位置。在这个程序中，我们不想并行启动多个块，而是想并行启动多个线程，因此使用 threadIdx.x 初始化 tid 为特定的线程 ID。这个程序的主函数如下所示：

int main(void) 
{
 //Defining Arrays for host
     float h_in[N], h_out[N];
     float *d_in, *d_out;
// allocate the memory on the cpu
     cudaMalloc((void**)&d_in, N * sizeof(float));
     cudaMalloc((void**)&d_out, N * sizeof(float));
 //Initializing Array
     for (int i = 0; i < N; i++) 
    {
         h_in[i] = i;
     }
 //Copy Array from host to device
     cudaMemcpy(d_in, h_in, N * sizeof(float), cudaMemcpyHostToDevice);
 //Calling square kernel with one block and N threads per block
     gpuSquare << <1, N >> >(d_in, d_out);
 //Coping result back to host from device memory
     cudaMemcpy(h_out, d_out, N * sizeof(float), cudaMemcpyDeviceToHost);
 //Printing result on console
     printf("Square of Number on GPU \n");
     for (int i = 0; i < N; i++) 
     {
         printf("The square of %f is %f\n", h_in[i], h_out[i]);
     }
 //Free up memory
     cudaFree(d_in);
     cudaFree(d_out);
     return 0;
 }

这个主函数遵循与向量加法程序相似的结构。你在这里会看到的一个区别是，我们正在并行地启动一个包含 N 个线程的单个块。程序输出如下：

每次你使用这种方式并行启动 N 个线程时，你应该注意每个块的最大线程数限制为 512 或 1,024。因此，N 的值应该小于这个值。如果你的设备每个块的最大线程数是 512，而 N 是 2,000，那么你不能写 << <1,2000 > >>。相反，你应该使用类似 << <4,500> >> 这样的东西。应该明智地选择块的数量和每个块中线程的数量。

总结一下，我们学习了如何处理向量，以及我们如何并行地启动多个块和多个线程。我们还看到，通过在 GPU 上执行向量操作，与在 CPU 上执行相同的操作相比，它提高了吞吐量。在本章的最后部分，我们将讨论线程并行执行时遵循的各种并行通信模式。

并行通信模式

当多个线程并行执行时，它们遵循一定的通信模式，这表明它们在哪里获取输入以及在内存中写入输出。我们将逐一讨论每种通信模式。这将帮助你识别与你的应用程序相关的通信模式以及如何编写相应的代码。

映射

在这种通信模式中，每个线程或任务取单个输入并产生单个输出。基本上，它是一个一对一的操作。前面章节中看到的向量加法程序和逐元素平方程序是映射模式的例子。映射模式的代码如下所示：

d_out[i] = d_in[i] * 2

Gather

在这种模式中，每个线程或任务有多个输入，并且它产生一个输出，该输出将被写入内存中的单个位置。假设你想编写一个程序来找到三个数的移动平均值；这是一个收集操作的例子。它从内存中获取三个输入，并将单个输出写入内存。因此，在输入端有数据重用。这基本上是一个多对一的操作。收集模式的代码如下所示：

out[i] = (in [i-1] + in[i] + in[i+1])/3

散列

在散列模式中，一个线程或任务取单个输入并计算在内存中应该写入输出的位置。数组排序是一个散列操作的例子。它也可以是一对多操作。散列模式的代码如下所示：

out[i-1] += 2 * in[i] and out[i+1] += 3*in[i]

模板

当线程或任务从一个数组的固定邻域集合中读取输入时，这被称为模板****通信模式。它在图像处理示例中非常有用，我们在 3x3 或 5x5 邻域窗口上工作。它是一种特殊的收集操作形式，因此代码语法与之相似。

转置

当输入以行主序矩阵的形式存在，而我们希望输出以列主序形式时，我们必须使用这种转置通信模式。如果你有一个数组结构并且想要将其转换为结构数组的形式，这尤其有用。它也是一种一对一操作。转置模式的代码如下所示：

out[i+j*128] = in [j +i*128]

在本节中，讨论了 CUDA 编程遵循的各种通信模式。找到与你的应用程序相关的通信模式并使用该模式的代码语法（如示例所示）是有用的。

摘要

总结来说，在本章中，你被介绍了 CUDA C 的编程概念以及如何使用 CUDA 进行并行计算。展示了 CUDA 程序可以高效且并行地运行在任何 NVIDIA GPU 硬件上。因此，CUDA 既高效又可扩展。详细讨论了在并行数据计算中需要的超出现有 ANSI C 函数的 CUDA API 函数。还通过一个简单的两个变量加法示例讨论了如何通过内核调用从主机代码调用设备代码、配置内核参数以及向内核传递参数。还展示了 CUDA 不保证块或线程的运行顺序以及哪个块被分配到哪个多处理器。此外，还讨论了利用 GPU 和 CUDA 的并行处理能力进行的向量操作。可以看出，通过在 GPU 上执行向量操作，与 CPU 相比，可以显著提高吞吐量。在最后一节中，详细讨论了并行编程中遵循的各种常见通信模式。然而，我们还没有讨论内存架构以及线程如何在 CUDA 中相互通信。如果一个线程需要其他线程的数据，那么可以做什么也没有讨论。因此，在下一章中，我们将详细讨论内存架构和线程同步。

问题

编写一个 CUDA 程序来减去两个数字。在内核函数中通过值传递参数。
编写一个 CUDA 程序来乘以两个数字。在内核函数中通过引用传递参数。
假设你想要并行启动 5,000 个线程。以三种不同的方式配置内核参数来完成此操作。每个块最多可以有 512 个线程。
对或错：程序员可以决定在设备上块将按何种顺序执行，以及块将被分配到哪个流多处理器？
编写一个 CUDA 程序以找出你的系统包含一个主次版本为 5.0 或更高版本的 GPU 设备。
编写一个 CUDA 程序来找到一个包含从 0 到 49 的数字的向量的立方。
对于以下应用，哪种通信模式是有用的？
1. 图像处理
2. 移动平均
3. 按升序排序数组
4. 在数组中查找数字的立方

第三章：线程、同步和内存

在上一章中，我们看到了如何编写 CUDA 程序，通过并行执行多个线程和块来利用 GPU 的处理能力。在所有程序中，直到上一章，所有线程都是相互独立的，并且多个线程之间没有通信。大多数现实生活中的应用程序需要中间线程之间的通信。因此，在本章中，我们将详细探讨如何在不同线程之间进行通信，并解释在处理相同数据的多线程之间的同步。我们将检查 CUDA 的分层内存架构以及如何使用不同的内存来加速 CUDA 程序。本章的最后部分解释了 CUDA 在向量点积和矩阵乘法中的一个非常有用应用，使用我们之前覆盖的所有概念。

本章将涵盖以下主题：

线程调用
CUDA 内存架构
全局、局部和缓存内存
共享内存和线程同步
原子操作
常量和纹理内存
点积和矩阵乘法示例

技术要求

本章要求熟悉基本的 C 或 C++编程语言以及前几章中解释的代码。本章中使用的所有代码都可以从以下 GitHub 链接下载：GitHub.com/PacktPublishing/Hands-On-GPU-Accelerated-Computer-Vision-with-OpenCV-and-CUDA。代码可以在任何操作系统上执行，尽管它只在 Windows 10 上进行了测试。

查看以下视频以查看代码的实际运行情况：

bit.ly/2prnGAD

线程

CUDA 在并行执行方面具有分层架构。内核执行可以在多个块上并行进行。每个块进一步分为多个线程。在上一章中，我们看到了 CUDA 运行时可以通过多次启动内核的相同副本来执行并行操作。我们看到了两种方法：要么并行启动多个块，每个块一个线程，要么启动单个块，并行启动多个线程。所以，你可能会有两个问题，我应该在我的代码中使用哪种方法？以及，并行启动的块和线程数量有什么限制？

这些问题的答案至关重要。正如我们将在本章后面看到的那样，同一块中的线程可以通过共享内存相互通信。因此，并行启动一个包含许多线程的块是有优势的，这样它们就可以相互通信。在上一章中，我们也看到了maxThreadPerBlock属性，它限制了每个块可以启动的线程数。对于最新的 GPU，其值为 512 或 1,024。同样，在第二种方法中，并行启动的最大块数限制为 65,535。

理想情况下，我们不是在每个单独的块中启动多个线程，或者不是在单个线程中启动多个块，而是在并行中启动多个块，每个块都有多个线程（可以等于maxThreadPerBlock）。所以，假设你想要在向量加法示例中并行启动 N = 50,000 个线程，这是我们上一章看到的。内核调用如下：

gpuAdd<< <((N +511)/512),512 > >>(d_a,d_b,d_c)

每个块的线程数最大为 512，因此总块数是通过将总线程数（N）除以 512 来计算的。但如果 N 不是 512 的准确倍数，那么 N 除以 512 可能会给出错误的块数，这个块数比实际数量少一个。因此，为了得到块数的下一个最高整数值，将 511 加到 N 上，然后再除以 512。这基本上是对除法进行向上取整操作。

现在，问题是，这对所有 N 的值都适用吗？答案是，很遗憾，不适用。从前面的讨论中可以看出，总块数不能超过 65,535。因此，在前面提到的内核调用中，如果(N+511)/512超过 65,535，那么代码将再次失败。为了克服这个问题，通过在内核代码中做一些修改，启动了少量块和线程，我们将在重写向量加法程序内核时进一步看到，如第二章中所述，使用 Cuda C 进行并行编程：

#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
//Defining number of elements in array
#define N 50000
__global__ void gpuAdd(int *d_a, int *d_b, int *d_c)
{
    //Getting index of current kernel
  int tid = threadIdx.x + blockIdx.x * blockDim.x; 

  while (tid < N)
    {
       d_c[tid] = d_a[tid] + d_b[tid];
       tid += blockDim.x * gridDim.x;
    }
}

这个内核代码与我们上一章中写的类似。它有两个修改。一个修改是在线程 ID 的计算中，第二个修改是在内核函数中包含while循环。线程 ID 计算的变化是由于并行启动多个线程和块。可以通过将块和线程视为一个二维矩阵来理解这个计算，其中块的数量等于行数，列数等于每个块的线程数。以下是一个例子，有三个块和三个线程/块，如下表所示：

我们可以通过使用blockIdx.x获取每个块的 ID，通过threadIdx.x命令获取当前块中每个线程的 ID。因此，对于显示为绿色的线程，块 ID 将是 2，线程 ID 将是 1。但如果我们想要一个在所有线程中唯一的索引呢？这可以通过将其块 ID 乘以每个块的总线程数（由blockDim.x给出）来计算，然后加上其线程 ID。这可以用以下数学公式表示：

tid = threadIdx.x + blockIdx.x * blockDim.x;

例如，在绿色部分，threadIdx.x = 1，blockIdx.x = 2，blockDim.x = 3 等于 tid = 7。这个计算非常重要，因为它将在你的代码中被广泛使用。

while循环被包含在代码中，因为当 N 非常大时，由于前面描述的限制，总线程数不能等于 N。因此，一个线程必须执行多个操作，这些操作由启动的总线程数分隔。这个值可以通过将blockDim.x乘以gridDim.x来计算，这分别给出了块和网格的维度。在while循环内部，线程 ID 通过这个偏移值增加。现在，这段代码将对任何 N 值都有效。为了完成程序，我们将为这段代码编写以下主函数：

int main(void) 
{
    //Declare host and device arrays
  int h_a[N], h_b[N], h_c[N];
  int *d_a, *d_b, *d_c;

    //Allocate Memory on Device
  cudaMalloc((void**)&d_a, N * sizeof(int));
  cudaMalloc((void**)&d_b, N * sizeof(int));
  cudaMalloc((void**)&d_c, N * sizeof(int));
    //Initialize host array
  for (int i = 0; i < N; i++) 
  {
    h_a[i] = 2 * i*i;
    h_b[i] = i;
  }

  cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice);
    //Kernel Call
  gpuAdd << <512, 512 >> >(d_a, d_b, d_c);

  cudaMemcpy(h_c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);
    //This ensures that kernel execution is finishes before going forward
  cudaDeviceSynchronize();
  int Correct = 1;
  printf("Vector addition on GPU \n");
  for (int i = 0; i < N; i++) 
  {
    if ((h_a[i] + h_b[i] != h_c[i]))
      { Correct = 0; }
  }
  if (Correct == 1)
  { 
    printf("GPU has computed Sum Correctly\n"); 
  }
  else
  { 
    printf("There is an Error in GPU Computation\n");
  }
    //Free up memory
  cudaFree(d_a);
  cudaFree(d_b);
   cudaFree(d_c);
  return 0;
}

再次强调，主要功能与我们上次写的内容非常相似。唯一的变化在于我们如何启动核函数。核函数以 512 个块的方式启动，每个块包含 512 个并行线程。这将解决 N 值较大的问题。我们不再打印一个非常长的向量的加法，而只打印一条指示计算结果是否正确的打印语句。代码的输出将如下所示：

本节解释了 CUDA 中的分层执行概念。下一节将进一步解释这个概念，通过解释分层内存架构。

内存架构

在 GPU 上执行代码被分配到流多处理器、块和线程中。GPU 有几个不同的内存空间，每个空间都有特定的特性和用途，以及不同的速度和范围。这个内存空间被分层划分为不同的部分，如全局内存、共享内存、局部内存、常量内存和纹理内存，并且它们可以从程序的不同点访问。这个内存架构在先前的图中显示：

如图中所示，每个线程都有自己的本地内存和寄存器文件。与处理器不同，GPU 核心拥有大量的寄存器来存储本地数据。当线程的数据不适合寄存器文件时，会使用本地内存。这两者都是每个线程独有的。寄存器文件是最快的内存。同一块中的线程共享内存，可以被该块中的所有线程访问。它用于线程间的通信。存在一个全局内存，可以被所有块和所有线程访问。全局内存具有较大的内存访问延迟。存在一种缓存的概念来加速这一操作。L1 和 L2 缓存如以下表格所示。存在一个只读的常量内存，用于存储常量和内核参数。最后，存在一个纹理内存，可以利用不同的二维或三维访问模式。

所有内存的特性总结在以下表格中：

内存	访问模式	速度	缓存？	作用域	生命周期
全局	读写	慢	是	主机和所有线程	整个程序
本地	读写	慢	是	每个线程	线程
寄存器	读写	快	-	每个线程	线程
共享	读写	快	否	每个块	块
常量	只读	慢	是	主机和所有线程	整个程序
纹理	只读	慢	是	主机和所有线程	整个程序

上述表格描述了所有内存的重要特性。作用域定义了程序可以使用此内存的部分，生命周期定义了该内存中的数据对程序可见的时间。除此之外，L1 和 L2 缓存也适用于 GPU 程序以实现更快的内存访问。

总结来说，所有线程都有一个寄存器文件，这是最快的。同一块中的多个线程有共享内存，比全局内存快。所有块都可以访问全局内存，这将是最慢的。常量和纹理内存用于特殊目的，将在下一节讨论。内存访问是程序快速执行中的最大瓶颈。

全局内存

所有块都可以读写全局内存。这种内存较慢，但可以从设备代码的任何地方访问。使用缓存的概念来加速对全局内存的访问。使用cudaMalloc分配的所有内存都将是一个全局内存。以下简单示例演示了您如何从程序中使用全局内存：

#include <stdio.h>
#define N 5

__global__ void gpu_global_memory(int *d_a)
{
  d_a[threadIdx.x] = threadIdx.x;
}

int main(int argc, char **argv)
{
  int h_a[N]; 
  int *d_a; 

  cudaMalloc((void **)&d_a, sizeof(int) *N);
  cudaMemcpy((void *)d_a, (void *)h_a, sizeof(int) *N, cudaMemcpyHostToDevice);

  gpu_global_memory << <1, N >> >(d_a); 
  cudaMemcpy((void *)h_a, (void *)d_a, sizeof(int) *N, cudaMemcpyDeviceToHost);

  printf("Array in Global Memory is: \n");
  for (int i = 0; i < N; i++) 
  {
    printf("At Index: %d --> %d \n", i, h_a[i]);
  }
  return 0;
}

此代码演示了您如何从设备代码写入全局内存。内存是通过主机代码中的cudaMalloc分配的，并将指向此数组的指针作为参数传递给内核函数。内核函数用线程 ID 的值填充这个内存块。然后将其复制回主机内存以打印。结果如下所示：

由于我们使用的是全局内存，这个操作将会较慢。有一些高级概念可以加快这个操作，稍后将会解释。在下一节中，我们将解释所有线程独有的局部内存和寄存器。

局部内存和寄存器

局部内存和寄存器文件是每个线程独有的。寄存器文件是每个线程可用的最快内存。当内核的变量不适合寄存器文件时，它们会使用局部内存。这被称为寄存器溢出。基本上，局部内存是全局内存的一部分，对每个线程来说是唯一的。与寄存器文件相比，访问局部内存会较慢。尽管局部内存被缓存在 L1 和 L2 缓存中，但寄存器溢出可能不会对程序产生负面影响。

下面是一个简单的程序，用于说明如何使用局部内存：

#include <stdio.h>
#define N 5

__global__ void gpu_local_memory(int d_in)
{
  int t_local; 
  t_local = d_in * threadIdx.x; 
  printf("Value of Local variable in current thread is: %d \n", t_local);
}
int main(int argc, char **argv)
{
  printf("Use of Local Memory on GPU:\n");
  gpu_local_memory << <1, N >> >(5); 
  cudaDeviceSynchronize();
  return 0;
}

t_local变量将属于每个线程，并存储在寄存器文件中。当这个变量在内核函数中进行计算时，计算将是最快的。前述代码的输出如下所示：

缓存内存

在最新的 GPU 上，每个多处理器都有一个 L1 缓存和一个 L2 缓存，这些缓存是所有多处理器共享的。全局和局部内存都使用这些缓存。由于 L1 缓存靠近线程执行，因此它非常快。如前所述的内存架构图所示，L1 缓存和共享内存使用相同的 64 KB。它们都可以配置为使用 64 KB 中的多少字节。所有全局内存访问都通过 L2 缓存进行。纹理内存和常量内存有自己的单独缓存。

线程同步

到目前为止，我们在这本书中看到的例子中，所有线程都是相互独立的。但在现实生活中，很少能找到线程在操作数据并终止时没有将结果传递给其他线程的例子。因此，线程之间必须有一些通信机制，这就是为什么在本节中解释了共享内存的概念。当许多线程并行工作并操作相同的数据或从相同的内存位置读取和写入时，所有线程之间必须进行同步。因此，本节还解释了线程同步。本节的最后部分解释了原子操作，这在读取-修改-写入条件下非常有用。

共享内存

共享内存位于芯片上，因此它比全局内存快得多。共享内存的延迟大约是全球未缓存内存延迟的 100 倍低。来自同一块的线程都可以访问共享内存。这在许多需要线程之间共享结果的程序中非常有用。然而，如果不进行同步，它也可能导致混乱或错误的结果。如果一个线程在另一个线程写入之前从内存中读取数据，可能会导致错误的结果。因此，内存访问应该得到适当的控制或管理。这是通过__syncthreads()指令完成的，它确保在程序前进之前所有对内存的write操作都已完成。这也被称为屏障。屏障的含义是所有线程都将到达这一行并等待其他线程完成。在所有线程都到达这个屏障之后，它们可以继续前进。为了演示共享内存和线程同步的使用，我们取了一个移动平均的例子。该内核函数如下所示：

#include <stdio.h>
__global__ void gpu_shared_memory(float *d_a)
{
  int i, index = threadIdx.x;
  float average, sum = 0.0f;
  //Defining shared memory
  __shared__ float sh_arr[10];

  sh_arr[index] = d_a[index];
 // This directive ensure all the writes to shared memory have completed

  __syncthreads();  
  for (i = 0; i<= index; i++) 
  { 
    sum += sh_arr[i]; 
  }
  average = sum / (index + 1.0f);
  d_a[index] = average;

    //This statement is redundant and will have no effect on overall code execution  
  sh_arr[index] = average;
}

移动平均操作不过是找到数组中所有元素的平均值，直到当前元素。许多线程将需要数组中的相同数据来进行计算。这是使用共享内存的理想情况，它将提供比全局内存更快的速度。这将减少每个线程的全局内存访问次数，从而降低程序的延迟。共享内存位置是通过__shared__指令定义的。在这个例子中，定义了十个浮点元素的共享内存。通常，共享内存的大小应该等于每个块中的线程数。在这里，我们正在处理一个包含 10 个元素的数组，因此我们采用了这个大小的共享内存。

下一步是将数据从全局内存复制到共享内存。所有线程将根据其线程 ID 索引的元素复制到共享数组。现在，这是一个共享内存写操作，在下一行中，我们将从这个共享数组中读取。因此，在继续之前，我们应该确保所有共享内存写操作都已完成。因此，让我们引入__synchronizethreads()屏障。

接下来，for循环使用共享内存中的值计算所有元素的平均值，并将结果存储在全局内存中，全局内存是通过当前线程 ID 索引的。最后一行也将计算出的值复制到共享内存中。这一行对代码的整体执行没有影响，因为共享内存的寿命直到当前块执行结束，这是块的最后一行。它只是用来演示这个关于共享内存的概念。现在，我们将尝试编写这个代码的主函数如下：

int main(int argc, char **argv)
{
   float h_a[10]; 
   float *d_a; 

      //Initialize host Array
   for (int i = 0; i < 10; i++) 
   {
     h_a[i] = i;
   }

    // allocate global memory on the device
    cudaMalloc((void **)&d_a, sizeof(float) * 10);

    // copy data from host memory  to device memory 
    cudaMemcpy((void *)d_a, (void *)h_a, sizeof(float) * 10,         cudaMemcpyHostToDevice);
    gpu_shared_memory << <1, 10 >> >(d_a);

    // copy the modified array back to the host
    cudaMemcpy((void *)h_a, (void *)d_a, sizeof(float) * 10, cudaMemcpyDeviceToHost);
    printf("Use of Shared Memory on GPU: \n");

    for (int i = 0; i < 10; i++) 
    {
      printf("The running average after %d element is %f \n", i, h_a[i]);
    }
    return 0;
}

在main函数中，在为宿主和设备数组分配内存之后，宿主数组被填充了从零到九的值。这些值被复制到设备内存中，在那里计算移动平均值，并将结果存储起来。设备内存中的结果被复制回宿主内存，然后打印到控制台。控制台上的输出如下所示：

本节演示了当多个线程使用同一内存位置的数据时共享内存的使用。下一节将演示使用atomic操作，这在读取-修改-写入操作中非常重要。

原子操作

考虑这样一种情况：大量线程试图修改内存的一小部分。这是一个经常发生的情况。当我们尝试执行读取-修改-写入操作时，这会引发更多的问题。这种操作的例子是d_out[i] ++,其中首先从内存中读取d_out[i]，然后增加，并写回内存。然而，当多个线程在相同的内存位置执行此操作时，可能会得到错误的结果。

假设一个内存位置的初始值为六，线程 p 和 q 都试图增加这个内存位置，那么最终的答案应该是八。但在执行时，可能会发生 p 和 q 线程同时读取这个值的情况，那么它们都会得到六这个值。它们将这个值增加到七，并将这个七存储在内存中。所以，最终答案不是八，而是七，这是错误的。这种错误可能带来的危险可以通过 ATM 取款的一个例子来理解。假设你在账户中有 5,000 卢比。你有两张相同的账户 ATM 卡。你和你的朋友同时去两个不同的 ATM 机取款 4,000 卢比。你们同时刷卡；所以，当 ATM 检查余额时，两个 ATM 都会显示 5,000 卢比的余额。当你们两人都取款 4,000 卢比时，那么两个机器都会查看初始余额，即 5,000 卢比。要取的金额小于余额，因此两个机器都会给出 4,000 卢比。尽管你的余额是 5,000 卢比，但你得到了 8,000 卢比，这是危险的。为了演示这种现象，我们取了一个大量线程尝试访问小数组的例子。这个例子的内核函数如下所示：

include <stdio.h>

#define NUM_THREADS 10000
#define SIZE 10

#define BLOCK_WIDTH 100

__global__ void gpu_increment_without_atomic(int *d_a)
{
  int tid = blockIdx.x * blockDim.x + threadIdx.x;

  // Each thread increment elements which wraps at SIZE
  tid = tid % SIZE;
  d_a[tid] += 1;
}

内核函数只是在d_a[tid] +=1行中增加内存位置。问题是这个内存位置增加了多少次。线程总数是 10,000，而数组的大小只有 10。我们通过将线程 ID 与数组大小进行取模操作来索引数组。因此，1,000 个线程将尝试增加相同的内存位置。理想情况下，数组的每个位置都应该增加 1,000 次。但正如我们将在输出中看到的，情况并非如此。在查看输出之前，我们将尝试编写main函数：

int main(int argc, char **argv)
{
  printf("%d total threads in %d blocks writing into %d array elements\n",
  NUM_THREADS, NUM_THREADS / BLOCK_WIDTH, SIZE);

  // declare and allocate host memory
  int h_a[SIZE];
  const int ARRAY_BYTES = SIZE * sizeof(int);
  // declare and allocate GPU memory
  int * d_a;
  cudaMalloc((void **)&d_a, ARRAY_BYTES);

  // Initialize GPU memory with zero value.
  cudaMemset((void *)d_a, 0, ARRAY_BYTES);
  gpu_increment_without_atomic << <NUM_THREADS / BLOCK_WIDTH, BLOCK_WIDTH >> >(d_a);

  // copy back the array of sums from GPU and print
  cudaMemcpy(h_a, d_a, ARRAY_BYTES, cudaMemcpyDeviceToHost);

  printf("Number of times a particular Array index has been incremented without atomic add is: \n");
  for (int i = 0; i < SIZE; i++)
  {
    printf("index: %d --> %d times\n ", i, h_a[i]);
  }
  cudaFree(d_a);
  return 0;
}

在main函数中，设备数组被声明并初始化为零。在这里，使用特殊的cudaMemSet函数在设备上初始化内存。这作为参数传递给内核，它增加这 10 个内存位置。在这里，总共启动了 10,000 个线程，分为 1,000 个块，每个块 100 个线程。内核执行后在设备上存储的答案被复制回主机，每个内存位置的价值在控制台上显示。

输出如下：

如前所述，理想情况下，每个内存位置应该增加 1,000 次，但大多数内存位置的价值为 16 和 17。这是因为许多线程同时读取相同的地址，因此增加相同的值并将其存储在内存中。由于线程执行的时机超出了程序员的控制，因此无法知道同时发生内存访问的次数。如果你再次运行你的程序，你的输出会与第一次相同吗？你的输出可能看起来像以下这样：

如你所猜，每次运行你的程序时，内存位置可能具有不同的值。这是因为设备上所有线程的随机执行导致的。

为了解决这个问题，CUDA 提供了一个名为atomicAdd操作的 API。这是一个阻塞操作，这意味着当多个线程试图访问相同的内存位置时，一次只能有一个线程可以访问该内存位置。其他线程必须等待这个线程完成并在内存上写入其答案。使用atomicAdd操作的内核函数如下所示：

#include <stdio.h>
#define NUM_THREADS 10000
#define SIZE 10
#define BLOCK_WIDTH 100

__global__ void gpu_increment_atomic(int *d_a)
{
  // Calculate thread index 
  int tid = blockIdx.x * blockDim.x + threadIdx.x;

  // Each thread increments elements which wraps at SIZE
  tid = tid % SIZE;
  atomicAdd(&d_a[tid], 1);
}

kernel函数与之前看到的非常相似。不是使用+=运算符增加内存位置，而是使用atomicAdd函数。它接受两个参数。第一个是我们想要增加的内存位置，第二个是这个位置需要增加的值。在这段代码中，1,000 个线程将再次尝试访问相同的位置；因此，当一个线程使用这个位置时，其他 999 个线程必须等待。这将增加执行时间方面的成本。使用atomic操作增加的main函数如下所示：

int main(int argc, char **argv)
{
  printf("%d total threads in %d blocks writing into %d array elements\n",NUM_THREADS, NUM_THREADS / BLOCK_WIDTH, SIZE);

  // declare and allocate host memory
  int h_a[SIZE];
  const int ARRAY_BYTES = SIZE * sizeof(int);

  // declare and allocate GPU memory
  int * d_a;
  cudaMalloc((void **)&d_a, ARRAY_BYTES);

   // Initialize GPU memory withzero value
  cudaMemset((void *)d_a, 0, ARRAY_BYTES);

  gpu_increment_atomic << <NUM_THREADS / BLOCK_WIDTH, BLOCK_WIDTH >> >(d_a);
    // copy back the array from GPU and print
  cudaMemcpy(h_a, d_a, ARRAY_BYTES, cudaMemcpyDeviceToHost);

  printf("Number of times a particular Array index has been incremented is: \n");
  for (int i = 0; i < SIZE; i++) 
  { 
     printf("index: %d --> %d times\n ", i, h_a[i]); 
  }

  cudaFree(d_a);
  return 0;
}

在main函数中，包含 10 个元素的数组被初始化为零值并传递给内核。但现在，内核将执行atomic add操作。因此，这个程序的输出应该是准确的。数组中的每个元素应该增加 1,000 次。以下将是输出结果：

如果你使用原子操作来测量程序的执行时间，它可能比使用全局内存的简单程序花费更长的时间。这是因为许多线程在原子操作中等待内存访问。使用共享内存可以帮助加快操作。此外，如果相同数量的线程访问更多的内存位置，那么原子操作将产生较少的时间开销，因为需要等待内存访问的线程数量更少。

在本节中，我们了解到原子操作有助于避免内存操作中的竞态条件，并使代码编写和理解更加简单。在下一节中，我们将解释两种特殊的内存类型，即常量和纹理内存，它们有助于加速某些类型的代码。

常量内存

CUDA 语言为程序员提供了一种另一种类型的内存，称为常量内存。NVIDIA 硬件提供了 64 KB 的这种常量内存，用于存储在整个内核执行过程中保持不变的数据。这种常量内存被缓存到芯片上，因此使用常量内存而不是全局内存可以加快执行速度。使用常量内存还将减少设备全局内存的带宽。在本节中，我们将了解如何在 CUDA 程序中使用常量内存。以执行简单数学运算a*x + b的简单程序为例，其中a和b是常量。该程序的kernel函数代码如下所示：

#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>

//Defining two constants
__constant__ int constant_f;
__constant__ int constant_g;
#define N 5

//Kernel function for using constant memory 
__global__ void gpu_constant_memory(float *d_in, float *d_out) 
{
  //Getting thread index for current kernel
  int tid = threadIdx.x; 
  d_out[tid] = constant_f*d_in[tid] + constant_g;
}

常量内存变量使用__constant__关键字定义。在前面的代码中，两个浮点变量constant_f和constant_g被定义为在整个内核执行过程中不会改变常量。第二点要注意的是，一旦变量被定义为常量，就不应该在内核函数中再次定义。内核函数使用这两个常量计算一个简单的数学运算。常量变量从main函数复制到内存中有一个特殊的方法。以下代码展示了这一点：

int main(void) 
{
  //Defining Arrays for host
  float h_in[N], h_out[N];
  //Defining Pointers for device
  float *d_in, *d_out;
  int h_f = 2;
  int h_g = 20;

  // allocate the memory on the cpu
  cudaMalloc((void**)&d_in, N * sizeof(float));
  cudaMalloc((void**)&d_out, N * sizeof(float));

  //Initializing Array
  for (int i = 0; i < N; i++) 
  {
    h_in[i] = i;
  }

  //Copy Array from host to device
  cudaMemcpy(d_in, h_in, N * sizeof(float), cudaMemcpyHostToDevice);
  //Copy constants to constant memory
  cudaMemcpyToSymbol(constant_f, &h_f, sizeof(int),0,cudaMemcpyHostToDevice);
  cudaMemcpyToSymbol(constant_g, &h_g, sizeof(int));

  //Calling kernel with one block and N threads per block
  gpu_constant_memory << <1, N >> >(d_in, d_out);

  //Coping result back to host from device memory
  cudaMemcpy(h_out, d_out, N * sizeof(float), cudaMemcpyDeviceToHost);

  //Printing result on console
  printf("Use of Constant memory on GPU \n");
  for (int i = 0; i < N; i++) 
  {
    printf("The expression for index %f is %f\n", h_in[i], h_out[i]);
  }

  cudaFree(d_in);
  cudaFree(d_out);
  return 0;
}

在main函数中，h_f和h_g常量在主机上定义并初始化，这些常量将被复制到常量内存中。使用cudaMemcpyToSymbol指令将这些常量复制到常量内存中以便内核执行。它有五个参数。第一个是目标，使用__constant__关键字定义。第二个是主机地址，第三个是传输的大小，第四个是内存偏移量，这里取为零，第五个是数据传输的方向，这里取为主机到设备。最后两个参数是可选的，因此在cudaMemcpyToSymbol指令的第二次调用中省略了它们。

代码的输出如下所示：

有一个需要注意的事项是常量内存是只读内存。这个例子只是用来解释从 CUDA 程序中使用常量内存。这不是常量内存的最佳使用。如前所述，常量内存有助于节省全局内存的内存带宽。要理解这一点，你必须理解 warp 的概念。一个 warp 是一组 32 个线程交织在一起并同步执行的集合。从常量内存的单次读取可以广播到半 warp，这可以减少多达 15 次内存事务。此外，常量内存被缓存，因此对附近位置的内存访问不会产生额外的内存事务。当每个包含 16 个线程的半 warp 在相同的内存位置上操作时，使用常量内存可以节省大量的执行时间。还应该注意的是，如果半 warp 线程使用完全不同的内存位置，那么使用常量内存可能会增加执行时间。因此，应该谨慎使用常量内存。

纹理内存

纹理内存是另一种只读内存，可以在以特定模式读取数据时加速程序并减少内存带宽。像常量内存一样，它也缓存于芯片上。这种内存最初是为渲染图形而设计的，但它也可以用于通用计算应用。当应用具有大量空间局部性的内存访问时，它非常有效。空间局部性的意义是每个线程很可能从其他附近线程读取的附近位置读取。这在图像处理应用中非常好，我们在其中处理 4 点连通性和 8 点连通性。线程通过访问内存位置进行二维空间局部性访问可能看起来像这样：

线程 0	线程 2
线程 1	线程 3

通用全局内存缓存将无法捕捉这种空间局部性，从而导致大量的内存流量到全局内存。纹理内存是为这种访问模式设计的，因此它只会从内存中读取一次，然后将其缓存起来，从而使执行速度大大加快。纹理内存支持一维和二维的fetch操作。在您的 CUDA 程序中使用纹理内存并不简单，尤其是对于那些不是编程专家的人来说。在本节中，解释了如何使用纹理内存复制数组值的简单示例。使用纹理内存的kernel函数解释如下：

#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>

#define NUM_THREADS 10
#define N 10

//Define texture reference for 1-d access
texture <float, 1, cudaReadModeElementType> textureRef;

__global__ void gpu_texture_memory(int n, float *d_out)
{
    int idx = blockIdx.x*blockDim.x + threadIdx.x;
    if (idx < n) {
      float temp = tex1D(textureRef, float(idx));
      d_out[idx] = temp;
    }
}

应该被读取的纹理内存部分由纹理引用定义。在代码中，它使用纹理 API 定义。它有三个参数。第一个参数指示纹理元素的数据类型。在这个例子中，它是一个float。第二个参数指示纹理引用的类型，可以是单维、二维等。在这里，它是一个单维引用。第三个参数指定读取模式，它是一个可选参数。请确保将此纹理引用声明为静态全局变量，并且它不应作为任何函数的参数传递。在内核函数中，存储在线程 ID 中的数据从这个纹理引用中读取，并复制到d_out全局内存指针。在这里，我们没有使用任何空间局部性，因为这个例子只是为了向您展示如何从 CUDA 程序中使用纹理内存。空间局部性将在下一章中解释，当我们看到一些使用 CUDA 的图像处理应用时。此例的main函数如下所示：

int main()
{
  //Calculate number of blocks to launch
  int num_blocks = N / NUM_THREADS + ((N % NUM_THREADS) ? 1 : 0);
  float *d_out;
  // allocate space on the device for the results
  cudaMalloc((void**)&d_out, sizeof(float) * N);
  // allocate space on the host for the results
  float *h_out = (float*)malloc(sizeof(float)*N);
  float h_in[N];
  for (int i = 0; i < N; i++) 
  {
    h_in[i] = float(i);
  }
  //Define CUDA Array
  cudaArray *cu_Array;
  cudaMallocArray(&cu_Array, &textureRef.channelDesc, N, 1);

  cudaMemcpyToArray(cu_Array, 0, 0, h_in, sizeof(float)*N, cudaMemcpyHostToDevice);

  // bind a texture to the CUDA array
  cudaBindTextureToArray(textureRef, cu_Array);

  gpu_texture_memory << <num_blocks, NUM_THREADS >> >(N, d_out);

  // copy result to host
  cudaMemcpy(h_out, d_out, sizeof(float)*N, cudaMemcpyDeviceToHost);
  printf("Use of Texture memory on GPU: \n");
  // Print the result
  for (int i = 0; i < N; i++) 
  {
    printf("Average between two nearest element is : %f\n", h_out[i]);
  }
  free(h_out);
  cudaFree(d_out);
  cudaFreeArray(cu_Array);
  cudaUnbindTexture(textureRef);
}

在main函数中，在声明和为主机和设备数组分配内存之后，主机数组使用从零到九的值进行初始化。在这个例子中，您将看到 CUDA 数组的首次使用。它们类似于普通数组，但它们是专门用于纹理的。它们对内核函数是只读的，可以通过使用cudaMemcpyToArray函数从主机写入设备内存，如前述代码所示。该函数中的第二个和第三个参数是宽度和高度偏移量，取值为 0，0，这意味着我们从左上角开始。它们是针对纹理内存读取优化的不透明内存布局。

cudaBindTextureToArray函数将纹理引用绑定到这个 CUDA 数组。这意味着，它从左上角开始将这个数组复制到纹理引用。绑定纹理引用后，调用内核，该内核使用这个纹理引用并计算要存储在设备内存上的数组。内核完成后，输出数组被复制回主机以在控制台上显示。当使用纹理内存时，我们必须使用cudaUnbindTexture函数从我们的代码中解除纹理的绑定。cudaFreeArray函数用于释放 CUDA 数组使用的内存。程序在控制台上显示的输出如下：

本节结束了我们对 CUDA 内存架构的讨论。当您根据您的应用程序合理地使用 CUDA 中可用的内存时，它可以极大地提高程序的性能。您需要仔细查看您应用程序中所有线程的内存访问模式，然后选择您应该为您的应用程序使用的内存。本章的最后一节简要描述了使用我们到目前为止所使用的所有概念的复杂 CUDA 程序。

点积和矩阵乘法示例

到目前为止，我们几乎已经学习了所有与 CUDA 基本并行编程相关的重要概念。在本节中，我们将向您展示如何编写 CUDA 程序来执行像点积和矩阵乘法这样的重要数学运算，这些运算几乎在所有应用中都会用到。这将利用我们之前看到的所有概念，并帮助您为您的应用程序编写代码。

点积

两个向量的点积是一个重要的数学运算。它还将解释 CUDA 编程中的一个重要概念，称为归约操作。两个向量的点积可以定义为如下：

(x1,x1,x3) . (y1,y2,y3) = x1y1 + x2y2 +x3y3

现在，如果您看到这个操作，它与向量上的逐元素加法操作非常相似。除了加法之外，您必须执行逐元素乘法。所有线程还必须继续运行它们所执行的乘法总和，因为所有单个乘法都需要相加以得到点积的最终答案。点积的答案将是一个单一的数字。在 CUDA 中，最终答案是原始两个数组的归约版本的操作称为归约操作。它在许多应用中非常有用。要在 CUDA 中执行此操作，我们将首先编写一个内核函数，如下所示：

#include <stdio.h>
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
#define N 1024
#define threadsPerBlock 512

__global__ void gpu_dot(float *d_a, float *d_b, float *d_c) 
{
  //Define Shared Memory
  __shared__ float partial_sum[threadsPerBlock];
  int tid = threadIdx.x + blockIdx.x * blockDim.x;
  int index = threadIdx.x;

  float sum = 0;
  while (tid < N) 
  {
    sum += d_a[tid] * d_b[tid];
    tid += blockDim.x * gridDim.x;
  }

  // set the partial sum in shared memory
  partial_sum[index] = sum;

  // synchronize threads in this block
  __syncthreads();

  //Calculate Patial sum for a current block using data in shared memory
  int i = blockDim.x / 2;
  while (i != 0) {
    if (index < i)
      {partial_sum[index] += partial_sum[index + i];}
    __syncthreads();
    i /= 2;
  }
  //Store result of partial sum for a block in global memory
  if (index == 0)
    d_c[blockIdx.x] = partial_sum[0];
}

kernel 函数接受两个输入数组作为输入，并将最终的局部和在第三个数组中存储。共享内存被定义为存储部分答案的中间答案。共享内存的大小等于每个块中的线程数，因为所有单独的块都将有这个共享内存的单独副本。之后，计算两个索引；第一个索引，用于计算唯一的线程 ID，类似于我们在向量加法示例中所做的。第二个索引用于在共享内存中存储部分乘积答案。同样，每个块都有一个单独的共享内存副本，所以只有用于索引共享内存的线程 ID 是给定块的。

while 循环将对由线程 ID 索引的元素执行逐元素乘法。它还将对偏移总线程数到当前线程 ID 的元素进行乘法。该元素的局部和存储在共享内存中。我们将使用这些来自共享内存的结果来计算单个块的局部和。因此，在读取这个共享内存块之前，我们必须确保所有线程都已经完成了对这个共享内存的写入。这通过使用 __syncthreads() 指令来确保。

现在，获取点积答案的一种方法是一个线程遍历所有这些部分和以获得最终答案。一个线程可以执行归约操作。这将需要 N 次操作来完成，其中 N 是要添加以获得最终答案的部分和的数量（等于每个块中的线程数）。

问题是，我们能否并行执行这个归约操作？答案是肯定的。想法是每个线程将添加两个部分和的元素并将答案存储在第一个元素的位置。由于每个线程结合了一个条目，所以操作可以在一半的条目中完成。现在，我们将重复这个操作，直到我们得到最终的答案，这个答案计算了整个块的局部和。这个操作的复杂度是 log2 ，这比一个线程执行归约操作的复杂度 N 要好得多。

解释的操作是通过以 while (i != 0) 开始的块来计算的。该块将当前线程的局部答案和偏移 blockdim/2 的线程的局部答案相加。它继续进行这种加法，直到我们得到一个最终的单一答案，这是给定块中所有部分乘积的总和。最终的答案存储在全局内存中。每个块都将有一个单独的答案存储在全局内存中，以便通过块 ID 索引，每个块都有一个唯一的块 ID。尽管如此，我们还没有得到最终的答案。这可以在 device 函数或 main 函数中执行。

通常，在归约操作的最后几个加法中需要的资源非常少。大部分 GPU 资源都处于空闲状态，这不是 GPU 的最佳使用。因此，单个块的各个部分的总和的最终加法操作是在main函数中完成的。main函数如下：

int main(void) 
{
  float *h_a, *h_b, h_c, *partial_sum;
  float *d_a, *d_b, *d_partial_sum;

  //Calculate number of blocks and number of threads
  int block_calc = (N + threadsPerBlock - 1) / threadsPerBlock;
  int blocksPerGrid = (32 < block_calc ? 32 : block_calc);
  // allocate memory on the cpu side
  h_a = (float*)malloc(N * sizeof(float));
  h_b = (float*)malloc(N * sizeof(float));
  partial_sum = (float*)malloc(blocksPerGrid * sizeof(float));

  // allocate the memory on the gpu
  cudaMalloc((void**)&d_a, N * sizeof(float));
  cudaMalloc((void**)&d_b, N * sizeof(float));
  cudaMalloc((void**)&d_partial_sum, blocksPerGrid * sizeof(float));

  // fill in the host mempory with data
  for (int i = 0; i<N; i++) {
    h_a[i] = i;
    h_b[i] = 2;
  }

  // copy the arrays to the device
  cudaMemcpy(d_a, h_a, N * sizeof(float), cudaMemcpyHostToDevice);
  cudaMemcpy(d_b, h_b, N * sizeof(float), cudaMemcpyHostToDevice);

  gpu_dot << <blocksPerGrid, threadsPerBlock >> >(d_a, d_b, d_partial_sum);

  // copy the array back to the host
  cudaMemcpy(partial_sum, d_partial_sum, blocksPerGrid * sizeof(float), cudaMemcpyDeviceToHost);

  // Calculate final dot prodcut
  h_c = 0;
  for (int i = 0; i<blocksPerGrid; i++) 
 {
    h_c += partial_sum[i];
  }

}

定义了三个数组，并为主机和设备分配了内存以存储输入和输出。两个主机数组在for循环内部初始化。一个数组初始化为从0到N，另一个数组初始化为常数2。计算网格中的块数和块中的线程数也已完成。这与我们在本章开头所做的是类似的。请注意，您也可以将这些值作为常数保留，就像我们在本章的第一个程序中所做的那样，以避免复杂性。

这些数组被复制到设备内存，并作为参数传递给kernel函数。kernel函数将返回一个数组，该数组包含由其块 ID 索引的各个块的乘积答案。这个数组被复制回主机到partial_sum数组中。点积的最终答案通过遍历这个partial_sum数组，使用从零开始的for循环到每个网格的块数来计算。最终的点积存储在h_c中。为了检查计算出的点积是否正确，可以在main函数中添加以下代码：

printf("The computed dot product is: %f\n", h_c);
#define cpu_sum(x) (x*(x+1))
  if (h_c == cpu_sum((float)(N - 1)))
  {
    printf("The dot product computed by GPU is correct\n");
  }
  else
  {
    printf("Error in dot product computation");
  }
  // free memory on the gpu side
  cudaFree(d_a);
  cudaFree(d_b);
  cudaFree(d_partial_sum);
  // free memory on the cpu side
  free(h_a);
  free(h_b);
  free(partial_sum);

答案通过数学计算的结果进行验证。在两个输入数组中，如果一个数组有从0到N-1的值，而第二个数组有一个常数值2，那么点积将是N*(N+1)。我们打印出数学计算出的点积答案，以及是否计算正确。最后释放主机和设备内存。程序的输出如下：

矩阵乘法

在使用 CUDA 在 GPU 上执行的第二重要的数学操作是矩阵乘法。当矩阵的大小非常大时，这是一个非常复杂的数学操作。应记住，对于矩阵乘法，第一个矩阵的列数应等于第二个矩阵的行数。矩阵乘法不是一个累积操作。为了避免复杂性，在这个例子中，我们取了一个相同大小的方阵。如果您熟悉矩阵乘法的数学，那么您可能会回忆起第一个矩阵的每一行将与第二个矩阵的所有列相乘。这将对第一个矩阵的所有行重复进行。如下所示：

同样的数据被多次重用，因此这是使用共享内存的理想情况。在本节中，我们将制作两个分别使用和不使用共享内存的单独的kernel函数。你可以比较两个内核的执行来了解共享内存如何提高程序的性能。我们首先从编写一个不使用共享内存的kernel函数开始：


#include <stdio.h>
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
#include <math.h>

//This defines size of a small square box or thread dimensions in one block
#define TILE_SIZE 2

//Matrix multiplication using non shared kernel
__global__ void gpu_Matrix_Mul_nonshared(float *d_a, float *d_b, float *d_c, const int size)
{
  int row, col;
  col = TILE_SIZE * blockIdx.x + threadIdx.x;
  row = TILE_SIZE * blockIdx.y + threadIdx.y;

  for (int k = 0; k< size; k++)
  {
    d_c[row*size + col] += d_a[row * size + k] * d_b[k * size + col];
  }
}

矩阵乘法使用二维线程执行。如果我们使用二维线程启动，每个线程执行输出矩阵的单个元素，那么最多可以乘以 16 x 16 的矩阵。如果大小大于这个值，那么计算将需要超过 512 个线程，这在大多数 GPU 上是不可能的。因此，我们需要启动多个块，每个块包含少于 512 个线程。为了实现这一点，输出矩阵被分成小正方形块，这两个方向上的维度都是TILE_SIZE。块中的每个线程将计算这个正方形块的元素。矩阵乘法的总块数将通过将矩阵的大小除以由TILE_SIZE定义的小正方形的大小来计算。

如果你理解了这一点，那么计算输出矩阵的行和列索引将会非常容易。这与我们到目前为止所做的是类似的，其中blockdim.x等于TILE_SIZE。现在，输出矩阵中的每个元素都将是一个矩阵第一行和一个矩阵第二列的点积。两个矩阵具有相同的大小，因此必须对等于大小变量的元素数量执行点积。因此，kernel函数中的for循环从0运行到size。

要计算两个矩阵的单独索引，考虑这个矩阵以行主序方式存储在系统内存中作为一个线性数组。这意味着第一行中的所有元素都放置在连续的内存位置，然后依次放置行，如下所示：

线性数组的索引可以通过其行 ID 乘以矩阵的大小加上其列 ID 来计算。因此，M[1,0]的索引将是 2，因为其行 ID 是 1，矩阵大小是 2，列 ID 是 0。这种方法用于计算两个矩阵中的元素索引。

要计算结果矩阵中[row, col]位置的元素，第一个矩阵中的索引将等于row*size + k，而对于第二个矩阵，它将是k*size + col。这是一个非常简单的kernel函数。在矩阵乘法中有大量的数据重用。这个函数没有利用共享内存的优势。因此，我们将尝试修改利用共享内存的kernel函数。修改后的kernel函数如下所示：

// shared
__global__ void gpu_Matrix_Mul_shared(float *d_a, float *d_b, float *d_c, const int size)
{
  int row, col;

  __shared__ float shared_a[TILE_SIZE][TILE_SIZE];

  __shared__ float shared_b[TILE_SIZE][TILE_SIZE];

  // calculate thread id
  col = TILE_SIZE * blockIdx.x + threadIdx.x;
  row = TILE_SIZE * blockIdx.y + threadIdx.y;

  for (int i = 0; i< size / TILE_SIZE; i++) 
  {
    shared_a[threadIdx.y][threadIdx.x] = d_a[row* size + (i*TILE_SIZE + threadIdx.x)];
    shared_b[threadIdx.y][threadIdx.x] = d_b[(i*TILE_SIZE + threadIdx.y) * size + col];
    }
    __syncthreads(); 

    for (int j = 0; j<TILE_SIZE; j++)
      d_c[row*size + col] += shared_a[threadIdx.x][j] * shared_b[j][threadIdx.y];
    __syncthreads(); // for synchronizing the threads

  }
}

定义了一个大小等于小方块块大小的两个共享内存，即TILE_SIZE，用于存储可重复使用的数据。行和列索引的计算方式与之前相同。首先，在第一个for循环中填充这个共享内存。之后，包含__syncthreads()，以确保只有当所有线程都完成写入后，才从共享内存中读取内存。最后一个for循环再次计算点积。由于这仅通过使用共享内存来完成，这大大减少了全局内存的内存流量，从而提高了程序在大矩阵维度上的性能。该程序的main函数如下所示：

int main()
{
   //Define size of the matrix
  const int size = 4;
   //Define host and device arrays
  float h_a[size][size], h_b[size][size],h_result[size][size];
  float *d_a, *d_b, *d_result; // device array
  //input in host array
  for (int i = 0; i<size; i++)
  {
    for (int j = 0; j<size; j++)
    {
      h_a[i][j] = i;
      h_b[i][j] = j;
    }
  }

  cudaMalloc((void **)&d_a, size*size*sizeof(int));
  cudaMalloc((void **)&d_b, size*size * sizeof(int));
  cudaMalloc((void **)&d_result, size*size* sizeof(int));
  //copy host array to device array
  cudaMemcpy(d_a, h_a, size*size* sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_b, h_b, size*size* sizeof(int), cudaMemcpyHostToDevice);
  //calling kernel
  dim3 dimGrid(size / TILE_SIZE, size / TILE_SIZE, 1);
  dim3 dimBlock(TILE_SIZE, TILE_SIZE, 1);

  gpu_Matrix_Mul_nonshared << <dimGrid, dimBlock >> > (d_a, d_b, d_result, size);
  //gpu_Matrix_Mul_shared << <dimGrid, dimBlock >> > (d_a, d_b, d_result, size);

  cudaMemcpy(h_result, d_result, size*size * sizeof(int), cudaMemcpyDeviceToHost);

  return 0;
}

在定义和分配主机和设备数组的内存之后，主机数组被填充了一些随机值。这些数组被复制到设备内存中，以便可以将其传递给kernel函数。使用dim3结构定义了网格块的数量和块线程的数量，其维度等于之前计算的值。您可以调用任何内核。将计算出的答案复制回主机内存。为了在控制台上显示输出，以下代码被添加到main函数中：

printf("The result of Matrix multiplication is: \n");

  for (int i = 0; i< size; i++)
  {
    for (int j = 0; j < size; j++)
    {
      printf("%f ", h_result[i][j]);
    }
    printf("\n");
  }
cudaFree(d_a)
cudaFree(d_b)
cudaFree(d_result)

用于在设备内存上存储矩阵的内存也被释放。控制台输出如下：

本节演示了在广泛应用的数学运算中使用的两个重要 CUDA 程序。它还解释了共享内存和多维线程的使用。

摘要

本章解释了多个块的启动，每个块都有来自内核函数的多个线程。它展示了选择大量线程的两个参数的方法。它还解释了 CUDA 程序可以使用的分层内存架构。最接近正在执行的线程的内存速度快，随着我们远离它，内存速度变慢。当多个线程想要相互通信时，CUDA 提供了使用共享内存的灵活性，使得同一块的线程可以相互通信。当多个线程使用相同的内存位置时，应该在内存访问之间进行同步；否则，最终结果将不会如预期。我们还看到了使用原子操作来完成这种同步的方法。如果某些参数在整个内核执行过程中保持不变，则可以将其存储在常量内存中以提高速度。当 CUDA 程序表现出某种通信模式，如空间局部性时，应使用纹理内存来提高程序的性能。总之，为了提高 CUDA 程序的性能，我们应该减少慢速内存的内存流量。如果这样做效率高，程序的性能可以得到显著提高。

在下一章中，我们将讨论 CUDA 流的概念，它与 CPU 程序中的多任务类似。我们还将讨论如何衡量 CUDA 程序的性能。它还将展示 CUDA 在简单的图像处理应用中的使用。

问题

假设你想要并行启动 100,000 个线程。在网格中块的数量和每个块中的线程数量最佳选择是什么，为什么？
编写一个 CUDA 程序，找出数组中每个元素的立方值，当数组中的元素数量为 100,000 时。
判断以下陈述是对还是错，并给出理由：局部变量之间的赋值运算符将比全局变量之间的赋值运算符更快。
注册溢出是什么？它如何损害你的 CUDA 程序的性能？
判断以下代码行是否会给出所需输出：d_out[i] = d_out[i-1]。
判断以下陈述是对还是错，并给出理由：原子操作会增加 CUDA 程序的执行时间。
在你的 CUDA 程序中使用纹理内存的理想通信模式有哪些？
在 if 语句中使用 __syncthreads 指令会有什么影响？

第四章：CUDA 的高级概念

在上一章中，我们探讨了 CUDA 的内存架构，并看到了如何有效地使用它来加速应用程序。到目前为止，我们还没有看到一种测量 CUDA 程序性能的方法。在本章中，我们将讨论如何使用 CUDA 事件来做到这一点。还将讨论 Nvidia Visual Profiler，以及如何在 CUDA 代码内部和调试工具中使用它来解决 CUDA 程序中的错误。我们还将讨论如何提高 CUDA 程序的性能。本章将描述如何使用 CUDA 流进行多任务处理，以及如何使用它们来加速应用程序。你还将学习如何使用 CUDA 加速数组排序算法。图像处理是一个需要在大约很短的时间内处理大量数据的领域，因此 CUDA 可以成为这类应用中操纵图像像素值的理想选择。本章描述了使用 CUDA 加速一个简单且广泛使用的图像处理函数——直方图计算。

本章将涵盖以下主题：

CUDA 中的性能测量
CUDA 中的错误处理
CUDA 程序的性能改进
CUDA 流及其如何用于加速应用程序
使用 CUDA 加速排序算法
使用 CUDA 介绍图像处理应用

技术要求

本章要求熟悉基本的 C 或 C++编程语言以及前几章中解释的所有代码示例。本章中使用的所有代码都可以从以下 GitHub 链接下载：github.com/PacktPublishing/Hands-On-GPU-Accelerated-Computer-Vision-with-OpenCV-and-CUDA。代码可以在任何操作系统上执行，尽管它只在 Windows 10 和 Ubuntu 上进行了测试。查看以下视频以查看代码的实际运行情况：

bit.ly/2Nt4DEy

CUDA 程序的性能测量

到目前为止，我们还没有明确确定 CUDA 程序的性能。在本节中，我们将看到如何使用 CUDA 事件来测量 CUDA 程序的性能，并使用 Nvidia Visual Profiler 来可视化性能。这在 CUDA 中是一个非常重要的概念，因为它将允许你从许多选项中选择特定应用程序的最佳性能算法。首先，我们将使用 CUDA 事件来测量性能。

CUDA 事件

我们可以使用 CPU 计时器来测量 CUDA 程序的性能，但它不会给出准确的结果。它将包括线程延迟开销和操作系统的调度，以及其他许多因素。使用 CPU 测量的时间也将取决于高精度 CPU 计时器的可用性。很多时候，当 GPU 内核运行时，主机正在执行异步计算，因此 CPU 计时器可能无法给出内核执行的正确时间。所以，为了测量 GPU 内核的计算时间，CUDA 提供了一个事件 API。

CUDA 事件是在您的 CUDA 程序中指定点记录的 GPU 时间戳。在这个 API 中，GPU 记录时间戳，消除了使用 CPU 计时器测量性能时存在的问题。使用 CUDA 事件测量时间有两个步骤：创建事件和记录事件。我们将记录两个事件，一个在代码的开始处，一个在结束处。然后，我们将尝试计算这两个事件之间时间差，这将给出代码的整体性能。

在您的 CUDA 代码中，您可以通过包含以下行来使用 CUDA 事件 API 来测量性能：

cudaEvent_t e_start, e_stop;
cudaEventCreate(&e_start);
cudaEventCreate(&e_stop);
cudaEventRecord(e_start, 0);
//All GPU code for which performance needs to be measured allocate the memory
cudaMalloc((void**)&d_a, N * sizeof(int));
cudaMalloc((void**)&d_b, N * sizeof(int));
cudaMalloc((void**)&d_c, N * sizeof(int));

  //Copy input arrays from host to device memory
cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice);

gpuAdd << <512, 512 >> >(d_a, d_b, d_c);
//Copy result back to host memory from device memory
cudaMemcpy(h_c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);
cudaDeviceSynchronize();
cudaEventRecord(e_stop, 0);
cudaEventSynchronize(e_stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, e_start, e_stop);
printf("Time to add %d numbers: %3.1f ms\n",N, elapsedTime);

我们将创建两个事件，e_start和e_stop，用于开始和结束代码。cudaEvent_t用于定义事件对象。要创建一个事件，我们将使用cudaEventCreate API。我们可以将事件对象作为参数传递给此 API。在代码的开始处，我们将记录 GPU 时间戳在e_start事件中；这将通过cudaEventRecord API 来完成。此函数的第二个参数是零，它表示 CUDA 流号，我们将在本章后面讨论。

在记录开始时的时间戳后，您可以开始编写您的 GPU 代码。在代码结束时，我们将在e_stop事件中再次记录时间。这将通过cudaEventRecord(e_stop, 0)行来完成。一旦我们记录了开始和结束时间，它们之间的差异应该会给我们代码的实际性能。但在这两个事件之间直接计算时间差仍然存在一个问题。

正如我们在前面的章节中讨论的那样，CUDA C 中的执行可以是异步的。当 GPU 执行内核时，CPU 可能会执行我们的代码的下一行，直到 GPU 完成其执行。所以，如果不同步 GPU 和 CPU 就直接测量时间可能会得到错误的结果。CudaEventRecord()会在其调用之前的所有 GPU 指令完成时记录一个时间戳。我们不应该在 GPU 上的先前工作完成之前读取e_stop事件。因此，为了同步 CPU 操作与 GPU，我们将使用cudaEventSynchronize(e_stop)。这确保了在e_stop事件中记录了正确的时间戳。

现在，为了计算这两个时间戳之间的差异，CUDA 提供了一个名为cudaEventElapsedTime的 API。它有三个参数。第一个是我们想要存储差异的变量，第二个是开始事件，第三个是结束事件。计算完这个时间后，我们将在下一行将其打印到控制台。我们将此性能测量代码添加到上一章中看到的向量加法代码中，使用了多个线程和块。添加这些行后的输出如下：

在 GPU 上添加 50,000 个元素所需的时间大约为 0.9 毫秒。此输出将取决于您的系统配置，因此您可能在红色框中得到不同的输出。因此，您可以将此性能测量代码包含在本书中看到的所有代码示例中，以测量它们的性能。您还可以通过使用此事件 API 来量化使用常量和纹理内存的性能提升。

应当记住，CUDA 事件只能用来测量设备代码块的执行时间。这仅包括内存分配、内存复制和内核执行。它不应用于测量主机代码的执行时间。因为 GPU 在事件 API 中记录时间，使用它来测量主机代码的性能可能会得到错误的结果。

Nvidia 视觉分析器

我们现在知道 CUDA 提供了一个有效的方法来提高并行计算应用程序的性能。然而，有时，即使将 CUDA 集成到您的应用程序中，代码的性能也可能不会提高。在这种情况下，可视化代码中哪个部分花费了最多时间完成是非常有用的。这被称为内核执行代码分析。Nvidia 提供了一个用于此的工具，并且它包含在标准的 CUDA 安装中。这个工具被称为Nvidia 视觉分析器。在 Windows 10 上的标准 CUDA 9.0 安装中，它可以在以下路径找到：C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0\libnvvp。您可以在该路径上运行nvvp应用程序，这将打开 Nvidia 视觉分析工具，如下所示：

此工具将执行您的代码，并根据您的 GPU 性能，为您提供每个内核的执行时间、代码中每个操作的详细时间戳、代码使用的内存以及内存带宽等详细信息。要为任何您开发的应用程序可视化和获取详细报告，您可以转到文件 -> 新会话。选择应用程序的.exe文件。我们选择了上一章中看到的向量加法示例。结果如下：

结果显示了程序中所有操作的计时。可以看到，cudaMalloc操作完成所需时间最长。它还显示了你的代码中每个操作执行的顺序。它显示内核只被调用了一次，平均需要 192.041 微秒来执行。内存复制操作的详细信息也可以可视化。从主机到设备的内存复制操作属性如下所示：

可以看到，当我们从主机复制两个数组到设备时，内存复制操作被调用了两次。总共复制的字节数为 400 KB，吞吐量为 1.693 GB/s。这个工具在内核执行分析中非常重要。它也可以用来比较两个内核的性能。它将显示导致你的代码性能下降的确切操作。

总结来说，在本节中，我们看到了两种测量和分析 CUDA 代码的方法。CUDA 事件是一个用于测量设备代码时序的效率 API。Nvidia Visual Profiler 提供了对 CUDA 代码的详细分析和性能分析，可用于性能分析。在下一节中，我们将看到如何处理 CUDA 代码中的错误。

CUDA 中的错误处理

我们还没有检查 CUDA 程序中 GPU 设备和内存的可用性。可能发生的情况是，当你运行 CUDA 程序时，GPU 设备不可用或内存不足。在这种情况下，你可能难以理解程序终止的原因。因此，在 CUDA 程序中添加错误处理代码是一个好习惯。在本节中，我们将尝试了解如何将此错误处理代码添加到 CUDA 函数中。当代码没有给出预期输出时，逐行检查代码的功能或通过在程序中添加断点来检查是有用的。这被称为调试。CUDA 提供了可以帮助的调试工具。因此，在接下来的部分，我们将看到 Nvidia 与 CUDA 一起提供的某些调试工具。

代码中的错误处理

当我们在第二章中讨论 CUDA API 函数时，使用 CUDA C 进行并行编程，我们看到了它们也返回一个标志，表示操作是否成功完成。这可以用来在 CUDA 程序中处理错误。当然，这不会帮助解决错误，但它会指示哪个 CUDA 操作导致了错误。将错误处理代码包装在 CUDA 函数中是一个非常良好的实践。以下是一个cudaMalloc函数的示例错误处理代码：

cudaError_t cudaStatus;
cudaStatus = cudaMalloc((void**)&d_a, sizeof(int));
if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
}

cudaError_t API 用于创建一个错误对象，该对象将存储所有 CUDA 操作的返回值。因此，cudaMalloc 函数的输出被分配给这个错误对象。如果错误对象不等于 cudaSuccess，则表示在设备上分配内存时出现了错误。这通过一个 if 语句来处理。它将在控制台上打印错误并跳转到程序的末尾。以下是一个在内存复制操作期间进行错误处理的包装代码示例：

cudaStatus = cudaMemcpy(d_a,&h_a, sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
  fprintf(stderr, "cudaMemcpy failed!");
  goto Error;
  }

再次强调，它与 cudaMalloc 的错误处理代码具有类似的结构。以下是一个内核调用包装代码的示例：

gpuAdd<<<1, 1>>>(d_a, d_b, d_c);
// Check for any errors launching the kernel
cudaStatus = cudaGetLastError();
if (cudaStatus != cudaSuccess) {
  fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
  goto Error;
}

内核调用不返回表示成功或失败的标志，因此它不会直接分配给错误对象。相反，如果在内核的启动过程中出现任何错误，则可以使用 cudaGetLastError() API 获取它，该 API 用于处理内核调用期间的错误。它被分配给 cudaStatus 错误对象，如果它不等于 cudaSuccess，它将在控制台上打印错误并跳转到程序的末尾。所有错误处理代码都会跳转到由 Error 标签定义的代码部分。它可以定义如下：

Error:
    cudaFree(d_a);

无论在程序中遇到任何错误，我们都会跳转到这个部分。我们将释放设备上分配的内存，然后退出 main 函数。这是一种编写 CUDA 程序的非常有效的方法。我们建议您使用这种方法来编写您的 CUDA 代码。之前没有解释这一点是为了避免在代码示例中引入不必要的复杂性。在 CUDA 程序中添加错误处理代码会使它们变得更长，但它能够确定是哪个 CUDA 操作在代码中引起问题。

调试工具

在编程中，我们可能会遇到两种类型的错误：语法错误和语义错误。语法错误可以通过编译器处理，但语义错误很难找到和调试。语义错误会导致程序出现意外的行为。当您的 CUDA 程序没有按预期工作，就需要逐行执行代码以可视化每行后的输出。这被称为调试。这对于任何类型的编程来说都是一个非常重要的操作。CUDA 提供了调试工具，有助于解决这类错误。

对于基于 Linux 的系统，Nvidia 提供了一个非常有用的调试器，称为 CUDA-GDB。它具有与用于 C 代码的正常 GDB 调试器类似的界面。它通过设置断点、检查 GPU 内存、检查块和线程等功能，帮助您在 GPU 上直接调试内核。它还提供了一个内存检查器来检查非法内存访问。

对于基于 Windows 的系统，Nvidia 提供了与 Microsoft Visual Studio 集成的 Nsight 调试器。同样，它具有在程序中添加断点和检查块或线程执行的功能。可以从 Visual Studio 内存界面查看设备的全局内存。

总结来说，在本节中，我们看到了两种处理 CUDA 中错误的方法。一种方法有助于解决与 GPU 硬件相关的错误，例如设备或内存不可用等 CUDA 程序中的错误。第二种使用调试的方法有助于当程序不符合预期时。在下一节中，我们将看到一些可以帮助提高 CUDA 程序性能的高级概念。

CUDA 程序的性能提升

在本节中，我们将看到一些基本指南，我们可以遵循这些指南来提高 CUDA 程序的性能。这些将逐一解释。

使用最佳数量的块和线程

我们在内核调用过程中看到了两个需要指定的参数：每个块的数量和每个块中的线程数。在内核调用期间，GPU 资源不应空闲；只有这样，它才能提供最佳性能。如果资源保持空闲，则可能会降低程序的性能。每个块和每个块中的线程数有助于保持 GPU 资源忙碌。研究表明，如果块的数量是 GPU 上多处理器数量的两倍，将提供最佳性能。GPU 上多处理器的总数可以通过使用设备属性找到，如第二章中所述，使用 CUDA C 进行并行编程。同样，每个块的最大线程数应等于maxThreadperblock设备属性。这些值仅作为指导。您可以通过调整这两个参数来获得应用程序中的最佳性能。

最大化算术效率

算术效率定义为数学运算次数与内存访问操作次数的比率。算术效率的值应尽可能高以获得良好的性能。可以通过最大化每个线程的运算次数和最小化每个线程在内存上的时间来实现。有时，最大化每个线程的运算次数的机会有限，但当然，您可以减少在内存上的时间。您可以通过将频繁访问的数据存储在快速内存中来最小化它。

我们在上一章中看到，局部内存和寄存器文件是 GPU 上可用的最快内存类型。因此，它们可以用来存储需要频繁访问的数据。我们还看到了使用共享内存、常量内存和纹理内存来提高性能。缓存也有助于减少内存访问时间。最终，如果我们减少全局内存的带宽，我们可以减少在内存上的时间。在提高 CUDA 程序性能方面，有效的内存使用非常重要，因为内存带宽是快速执行中的最大瓶颈。

使用归一化或步进内存访问

合并内存访问意味着每个线程都读取或写入连续的内存位置。当使用这种内存访问方法时，GPU 效率最高。如果线程使用偏移量为常数的内存位置，则这被称为步进内存访问。它仍然比随机内存访问有更好的性能。因此，如果您在程序中尝试使用合并内存访问，它可以显著提高性能。以下是一些这些内存访问模式的示例：

Coalesce Memory Access: d_a[i] = a
Strided Memory Access: d_a[i*2] = a

避免线程发散

当内核中的所有线程调用不同的执行路径时，会发生线程发散。它可以在以下内核代码场景中发生：

Thread divergence by way of branching
tid = ThreadId
if (tid%2 == 0)
{ 
  Some Branch code;
}
else
{
  Some other code; 
}
Thread divergence by way of looping 
Pre-loop code
for (i=0; i<tid;i++)
{
  Some loop code;
}
Post loop code;

在第一个代码片段中，由于if语句中的条件，存在针对奇数和偶数线程的单独代码。这使得奇数和偶数线程遵循不同的执行路径。在if语句之后，这些线程将再次合并。这将产生时间开销，因为快速线程将不得不等待慢速线程。

在第二个示例中，使用for循环，每个线程运行for循环的迭代次数不同，因此所有线程完成所需的时间不同。循环后的代码必须等待所有这些线程完成。这将产生时间开销。因此，尽可能避免在您的代码中这种类型的线程发散。

使用页面锁定主机内存

在此之前的每个示例中，我们使用malloc函数在主机上分配内存，这在主机上分配标准可分页内存。CUDA 提供了一个名为cudaHostAlloc()的另一个 API，它分配页面锁定主机内存或有时称为固定内存。它保证操作系统永远不会将此内存从磁盘页出，并且它将保留在物理内存中。因此，任何应用程序都可以访问缓冲区的物理地址。这种属性有助于 GPU 通过直接内存访问（DMA）将数据从主机复制到主机，而无需 CPU 干预。这有助于提高内存传输操作的性能。但是，应该小心使用固定内存，因为这种内存不会被换出到磁盘；您的系统可能耗尽内存。它可能影响系统上运行的其他应用程序的性能。您可以使用此 API 分配用于通过Memcpy操作将数据传输到设备的内存。使用此 API 的语法如下：

Allocate Memory: cudaHostAlloc ( (void **) &h_a, sizeof(*h_a), cudaHostAllocDefault);
Free Memory: cudaFreeHost(h_a);

cudaHostAlloc的语法类似于简单的malloc函数。最后一个参数，cudaHostAllocDefault，是一个用于修改固定内存行为的标志。cudaFreeHost用于释放使用cudaHostAlloc函数分配的内存。

CUDA 流

我们已经看到，当单个指令对多个数据项进行操作时，GPU 在数据并行性方面提供了极大的性能提升。我们还没有看到任务并行性，其中多个相互独立的内核函数并行运行。例如，一个函数可能正在计算像素值，而另一个函数正在从互联网上下载某些内容。我们知道 CPU 提供了非常灵活的方法来处理这种类型的任务并行性。GPU 也提供了这种能力，但它的灵活性不如 CPU。这种任务并行性是通过使用 CUDA 流实现的，我们将在本节中详细探讨。

CUDA 流实际上是一个 GPU 操作的队列，这些操作按特定顺序执行。这些函数包括内核函数、内存复制操作和 CUDA 事件操作。它们被添加到队列中的顺序将决定它们的执行顺序。每个 CUDA 流可以被视为一个单独的任务，因此我们可以启动多个流以并行执行多个任务。我们将在下一节中查看 CUDA 中多个流的工作方式。

使用多个 CUDA 流

我们将通过在上一章开发的向量加法程序中使用多个 CUDA 流来理解 CUDA 流的工作原理。这个内核函数如下所示：

#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
//Defining number of elements in Array
#define N 50000

//Defining Kernel function for vector addition
__global__ void gpuAdd(int *d_a, int *d_b, int *d_c) {
  //Getting block index of current kernel

  int tid = threadIdx.x + blockIdx.x * blockDim.x;
  while (tid < N)
  {
    d_c[tid] = d_a[tid] + d_b[tid];
    tid += blockDim.x * gridDim.x;
  }
}

内核函数与我们之前开发的类似。它只是多个流将并行执行这个内核。需要注意的是，并非所有 GPU 设备都支持 CUDA 流。支持deviceOverlap属性的 GPU 设备可以同时执行内存传输操作和内核执行。这个属性将在 CUDA 流中用于任务并行。在继续此代码之前，请确保您的 GPU 设备支持此属性。您可以使用第二章中的代码，使用 CUDA C 进行并行编程，来验证此属性。我们将使用两个并行流，它们将并行执行此内核，并对输入数据的一半进行操作。我们将在主函数中首先创建这两个流，如下所示：

int main(void) {
  //Defining host arrays
  int *h_a, *h_b, *h_c;
  //Defining device pointers for stream 0
  int *d_a0, *d_b0, *d_c0;
  //Defining device pointers for stream 1
 int *d_a1, *d_b1, *d_c1;
 cudaStream_t stream0, stream1;
 cudaStreamCreate(&stream0);
 cudaStreamCreate(&stream1);

cudaEvent_t e_start, e_stop;
 cudaEventCreate(&e_start);
  cudaEventCreate(&e_stop);
  cudaEventRecord(e_start, 0);

使用cudaStream_t和cudaStreamCreate API 定义了两个流对象，stream 0和stream 1。我们还定义了主机指针和两组设备指针，它们将分别用于每个流。我们定义并创建了两个事件来测量此程序的性能。现在，我们需要为这些指针分配内存。代码如下：

  //Allocate memory for host pointers
  cudaHostAlloc((void**)&h_a, 2*N* sizeof(int),cudaHostAllocDefault);
 cudaHostAlloc((void**)&h_b, 2*N* sizeof(int), cudaHostAllocDefault);
 cudaHostAlloc((void**)&h_c, 2*N* sizeof(int), cudaHostAllocDefault);
  //Allocate memory for device pointers
  cudaMalloc((void**)&d_a0, N * sizeof(int));
  cudaMalloc((void**)&d_b0, N * sizeof(int));
  cudaMalloc((void**)&d_c0, N * sizeof(int));
  cudaMalloc((void**)&d_a1, N * sizeof(int));
  cudaMalloc((void**)&d_b1, N * sizeof(int));
  cudaMalloc((void**)&d_c1, N * sizeof(int));
  for (int i = 0; i < N*2; i++) {
    h_a[i] = 2 * i*i;
    h_b[i] = i;
  }

CUDA 流在进行内存复制操作时需要访问页锁定内存，因此我们使用cudaHostAlloc函数而不是简单的malloc来定义主机内存。我们在上一节中看到了页锁定内存的优势。使用cudaMalloc分配了两组设备指针的内存。需要注意的是，主机指针持有全部数据，因此其大小为2*N*sizeof(int)，而每个设备指针只操作一半的数据元素，因此其大小仅为N*sizeof(int)。我们还用一些随机值初始化了主机数组以进行加法操作。现在，我们将尝试在两个流中同时排队内存复制操作和内核执行操作。相应的代码如下：

//Asynchrnous Memory Copy Operation for both streams
cudaMemcpyAsync(d_a0, h_a , N * sizeof(int), cudaMemcpyHostToDevice, stream0);
cudaMemcpyAsync(d_a1, h_a+ N, N * sizeof(int), cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d_b0, h_b , N * sizeof(int), cudaMemcpyHostToDevice, stream0);
cudaMemcpyAsync(d_b1, h_b + N, N * sizeof(int), cudaMemcpyHostToDevice, stream1);

//Kernel Call     
gpuAdd << <512, 512, 0, stream0 >> > (d_a0, d_b0, d_c0);
gpuAdd << <512, 512, 0, stream1 >> > (d_a1, d_b1, d_c1);

//Copy result back to host memory from device memory
cudaMemcpyAsync(h_c , d_c0, N * sizeof(int), cudaMemcpyDeviceToHost, stream0);
cudaMemcpyAsync(h_c + N, d_c1, N * sizeof(int), cudaMemcpyDeviceToHost, stream0);

我们不是使用简单的cudaMemcpy API，而是使用cudaMemcpyAsync API，它用于异步内存传输。它将一个内存复制操作的请求排队到由函数的最后一个参数指定的给定流中。当这个函数返回时，内存复制操作可能还没有开始，因此它被称为异步操作。它只是将内存复制的请求放入队列中。正如我们可以在内存复制操作中看到的那样，stream0操作从0到N的数据，而stream 1操作从N+1到2N的数据。

在流操作中，操作顺序很重要，因为我们希望内存复制操作与内核执行操作重叠。因此，我们不是先排队所有stream0操作，然后排队stream 1操作，而是首先在两个流中排队内存复制操作，然后排队内核计算操作。这将确保内存复制和内核计算相互重叠。如果这两个操作花费相同的时间，我们可以实现两倍的速度提升。我们可以通过查看以下图表来更好地了解操作顺序：

时间从上到下增加。我们可以看到，在同一时间段内执行了两个内存复制操作和内核执行操作，这将加速你的程序。我们还看到，由cudaMemcpyAsync定义的内存复制操作是异步的；因此，当一个流返回时，内存复制操作可能还没有开始。如果我们想使用最后一个内存复制操作的结果，那么我们必须等待两个流完成它们的队列操作。这可以通过使用以下代码来确保：

cudaDeviceSynchronize();
cudaStreamSynchronize(stream0);
cudaStreamSynchronize(stream1);

cudaStreamSynchronize确保在继续到下一行之前，流中的所有操作都已完成。为了测量代码的性能，我们插入以下代码：

cudaEventRecord(e_stop, 0);
cudaEventSynchronize(e_stop);
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, e_start, e_stop);
printf("Time to add %d numbers: %3.1f ms\n",2* N, elapsedTime);

它将记录停止时间，并根据开始和停止时间之间的差异，计算该程序的总体执行时间，并在控制台上打印输出。为了检查程序是否计算了正确的输出，我们将插入以下代码进行验证：

int Correct = 1;
printf("Vector addition on GPU \n");
//Printing result on console
for (int i = 0; i < 2*N; i++) 
{
  if ((h_a[i] + h_b[i] != h_c[i]))
  {
    Correct = 0;
  }
}

if (Correct == 1)
{
  printf("GPU has computed Sum Correctly\n");
}
else
{
  printf("There is an Error in GPU Computation\n");
}
//Free up memory
cudaFree(d_a0);
cudaFree(d_b0);
cudaFree(d_c0);
cudaFree(d_a0);
cudaFree(d_b0);
cudaFree(d_c0);
cudaFreeHost(h_a);
cudaFreeHost(h_b);
cudaFreeHost(h_c);
return 0;
}

验证代码与我们之前看到的类似。使用cudaFree释放设备上分配的内存，使用cudaHostAlloc在主机上分配的内存使用cudaFreeHost函数释放。这是强制性的，否则您的系统可能会很快耗尽内存。程序输出如下所示：

如前一个截图所示，需要 0.9 毫秒来添加 100,000 个元素，这是在没有流的情况下代码的两倍增加，如本章第一部分所示，添加 50,000 个数字需要 0.9 毫秒。

总结来说，在本节中我们看到了 CUDA 流，它有助于在 GPU 上实现任务并行。在流中排队操作的顺序对于使用 CUDA 流实现加速非常重要。

使用 CUDA 加速排序算法

排序算法在许多计算应用中被广泛使用。有许多排序算法，例如枚举或排名排序、冒泡排序和归并排序。所有算法都有不同的复杂度级别，因此对给定数组进行排序所需的时间不同。对于大型数组，所有算法都需要很长时间才能完成。如果可以使用 CUDA 进行加速，那么它将对任何计算应用都有很大帮助。

为了展示 CUDA 如何加速不同的排序算法，我们将实现一个排名排序算法。

枚举或排名排序算法

在这个算法中，我们计算数组中的每个元素，以找出数组中有多少元素小于当前元素。从那里，我们可以得到当前元素在排序数组中的位置。然后，我们将此元素放在那个位置。我们重复这个过程，直到数组中的所有元素，以得到一个排序数组。这被实现为kernel函数，如下所示：

#include "device_launch_parameters.h"
#include <stdio.h>

#define arraySize 5
#define threadPerBlock 5
//Kernel Function for Rank sort
__global__ void addKernel(int *d_a, int *d_b)
{
  int count = 0;
  int tid = threadIdx.x;
  int ttid = blockIdx.x * threadPerBlock + tid;
  int val = d_a[ttid];
  __shared__ int cache[threadPerBlock];
  for (int i = tid; i < arraySize; i += threadPerBlock) {
    cache[tid] = d_a[i];
    __syncthreads();
    for (int j = 0; j < threadPerBlock; ++j)
      if (val > cache[j])
        count++;
        __syncthreads();
  }
  d_b[count] = val;
}

Kernel函数接受两个数组作为参数。d_a是输入数组，d_b是输出数组。count变量被取用，它存储当前元素在排序数组中的位置。当前线程在块中的索引存储在tid中，所有块中唯一的线程索引存储在ttid中。使用共享内存来减少从全局内存访问数据的时间。共享内存的大小等于块中线程的数量，如前所述。value变量持有当前元素。共享内存被填充为全局内存中的值。这些值与value变量进行比较，并将小于的值的数量存储在count变量中。这会一直持续到数组中的所有元素都与value变量进行比较。循环结束后，count变量有元素在排序数组中的位置，并将当前元素存储在输出数组d_b中的那个位置。

该代码的main函数如下：

int main()
{
    //Define Host and Device Array
  int h_a[arraySize] = { 5, 9, 3, 4, 8 };
  int h_b[arraySize];
  int *d_a, *d_b;

    //Allocate Memory on the device 
  cudaMalloc((void**)&d_b, arraySize * sizeof(int));
  cudaMalloc((void**)&d_a, arraySize * sizeof(int));

    // Copy input vector from host memory to device memory.
  cudaMemcpy(d_a, h_a, arraySize * sizeof(int), cudaMemcpyHostToDevice);

    // Launch a kernel on the GPU with one thread for each element.
  addKernel<<<arraySize/threadPerBlock, threadPerBlock>>>(d_a, d_b);

    //Wait for device to finish operations
  cudaDeviceSynchronize();
    // Copy output vector from GPU buffer to host memory.
  cudaMemcpy(h_b, d_b, arraySize * sizeof(int), cudaMemcpyDeviceToHost);
  printf("The Enumeration sorted Array is: \n");
  for (int i = 0; i < arraySize; i++) 
  {
    printf("%d\n", h_b[i]);
  }
    //Free up device memory
  cudaFree(d_a);
  cudaFree(d_b);
  return 0;
}

main 函数你现在应该已经很熟悉了。我们正在定义主机和设备数组，并在设备上为设备数组分配内存。主机数组使用一些随机值初始化，并将其复制到设备的内存中。通过传递设备指针作为参数来启动内核。内核通过排名排序算法计算排序后的数组，并将其返回到主机。这个排序后的数组如下所示在控制台上打印：

这是一个非常简单的情况，你可能看不到 CPU 和 GPU 之间有任何性能提升。但是，如果你继续增加 arraySize 的值，那么 GPU 将会极大地提高这个算法的性能。对于大小等于 15,000 的数组，它可以实现百倍的性能提升。

排名排序是可用的最简单的排序算法。这次讨论将帮助你开发其他排序算法的代码，例如冒泡排序和归并排序。

使用 CUDA 进行图像处理

现在，我们生活在一个高清摄像头传感器时代，它可以捕捉高分辨率的图像。一个图像可以达到 1920 x 1920 像素的大小。因此，在计算机上实时处理这些像素需要每秒执行数十亿次的浮点运算。即使是速度最快的 CPU 也难以做到这一点。GPU 可以在这种情况下提供帮助。它提供了高计算能力，这可以通过 CUDA 在你的代码中利用。

在计算机中，图像以多维数组的形式存储，灰度图像有两个维度，彩色图像有三个维度。CUDA 也支持多维网格块和线程。因此，我们可以通过启动多维块和线程来处理图像，就像之前看到的那样。块和线程的数量可以取决于图像的大小。它也将取决于你的 GPU 规格。如果它支持每个块 1,024 个线程，那么可以启动每个块 32 x 32 个线程。块的数量可以通过将图像大小除以这些线程的数量来确定。正如之前多次讨论的那样，参数的选择会影响你代码的性能。因此，它们应该被适当地选择。

将用 C 或 C++ 开发的简单图像处理代码转换为 CUDA 代码非常容易。即使是不经验验的程序员也可以通过遵循一个固定的模式来完成。图像处理代码有一个固定的模式，如下面的代码所示：

for (int i=0; i < image_height; i++)
{
   for (int j=0; j < image_width; j++)
   {
      //Pixel Processing code for pixel located at (i,j)
   }
}

图像不过是存储在计算机上的多维矩阵，因此从图像中获取单个像素值需要使用嵌套的 for 循环来遍历所有像素。为了将此代码转换为 CUDA，我们希望启动与图像中像素数量相等的线程数量。在 kernel 函数中，可以通过以下代码在线程中获取像素值：

int i = blockIdx.y * blockDim.y + threadIdx.y;
int j = blockIdx.x * blockDim.x + threadIdx.x;

i和j值可以用作图像数组的索引来查找像素值。所以，如前述代码所示，通过将for循环转换为线程索引的简单转换过程，我们可以编写 CUDA 程序的设备代码。从下一节开始，我们将使用OpenCV库开发许多图像处理应用程序。在本章中，我们不会涵盖实际的图像操作，但我们将通过开发一个用于计算直方图这一重要统计操作的 CUDA 程序来结束本章。直方图计算对于图像处理应用程序也非常重要。

使用 CUDA 在 GPU 上计算直方图

直方图是一个非常重要的统计概念，在机器学习、计算机视觉、数据科学和图像处理等多种应用中使用。它表示给定数据集中每个元素频率的计数。它显示了哪些数据项出现频率最高，哪些出现频率最低。你也可以通过查看直方图的值来了解数据的分布。在本节中，我们将开发一个算法，用于计算给定数据分布的直方图。

我们将首先在 CPU 上计算直方图，这样你可以了解如何计算直方图。假设我们有一组包含 1,000 个元素的数，每个元素的价值在 0 到 15 之间。我们想要计算这个分布的直方图。在 CPU 上计算这个计算的示例代码如下：

int h_a[1000] = Random values between 0 and 15

int histogram[16];
for (int i = 0; i<16; i++)
{ 
   histogram[i] = 0;
}
for (i=0; i < 1000; i++)
{
   histogram[h_a[i]] +=1;
}

我们有 1,000 个数据元素，它们存储在h_a中。h_a数组包含0到15之间的值；它有 16 个不同的值。因此，bin 的数量，即需要计算直方图的唯一值的数量，是 16。因此，我们定义了一个大小等于 bin 数量的直方图数组，用于存储最终的直方图。这个数组需要初始化为零，因为它将在每次发生时递增。这是在从0到 bin 数量的第一个for循环中完成的。

对于直方图的计算，我们需要遍历h_a中的所有元素。在h_a中找到的任何值，都需要增加该直方图数组中特定索引的值。这是通过第二个for循环完成的，该循环从0到数组大小运行，并增加由h_a中找到的值索引的直方图数组。在for循环完成后，直方图数组将包含0到15之间每个元素的频率。

现在，我们将为 GPU 开发相同的代码。我们将尝试使用三种不同的方法来开发这个代码。前两种方法的内核代码如下：

#include <stdio.h>
#include <cuda_runtime.h>

#define SIZE 1000
#define NUM_BIN 16

__global__ void histogram_without_atomic(int *d_b, int *d_a)
{
  int tid = threadIdx.x + blockDim.x * blockIdx.x;
  int item = d_a[tid];
  if (tid < SIZE)
  {
    d_b[item]++;
  }
 }

__global__ void histogram_atomic(int *d_b, int *d_a)
{
  int tid = threadIdx.x + blockDim.x * blockIdx.x;
  int item = d_a[tid];
  if (tid < SIZE)
  {
    atomicAdd(&(d_b[item]), 1);
  }
}

第一个函数是直方图计算的最简单内核函数。每个线程都在操作一个数据元素。使用线程 ID 作为索引从输入数组中获取数据元素的值。这个值被用作d_b输出数组的索引，该数组被递增。d_b数组应该包含输入数据中每个值（0到15）的频率。但如果你回想一下第三章，线程、同步和内存，这可能不会给你一个正确答案，因为许多线程正在同时尝试修改相同的内存位置。在这个例子中，1,000 个线程正在同时尝试修改 16 个内存位置。我们需要在这种情况下使用原子的add操作。

第二个设备函数是使用原子add操作开发的。这个内核函数将给出正确答案，但完成所需的时间会更长，因为原子操作是一个阻塞操作。当有一个线程正在使用特定的内存位置时，所有其他线程都必须等待。因此，这个第二个内核函数将增加开销时间，使其比 CPU 版本还要慢。为了完成代码，我们将尝试按照以下方式编写它的main函数：

int main()
{

  int h_a[SIZE];
  for (int i = 0; i < SIZE; i++) {

  h_a[i] = i % NUM_BIN;
  }
  int h_b[NUM_BIN];
  for (int i = 0; i < NUM_BIN; i++) {
    h_b[i] = 0;
  }

  // declare GPU memory pointers
  int * d_a;
  int * d_b;

  // allocate GPU memory
  cudaMalloc((void **)&d_a, SIZE * sizeof(int));
  cudaMalloc((void **)&d_b, NUM_BIN * sizeof(int));

  // transfer the arrays to the GPU
  cudaMemcpy(d_a, h_a, SIZE * sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_b, h_b, NUM_BIN * sizeof(int), cudaMemcpyHostToDevice);

  // launch the kernel

  //histogram_without_atomic << <((SIZE+NUM_BIN-1) / NUM_BIN), NUM_BIN >> >(d_b, d_a);
  histogram_atomic << <((SIZE+NUM_BIN-1) / NUM_BIN), NUM_BIN >> >(d_b, d_a);

  // copy back the sum from GPU
  cudaMemcpy(h_b, d_b, NUM_BIN * sizeof(int), cudaMemcpyDeviceToHost);
  printf("Histogram using 16 bin without shared Memory is: \n");
  for (int i = 0; i < NUM_BIN; i++) {
    printf("bin %d: count %d\n", i, h_b[i]);
  }

  // free GPU memory allocation
  cudaFree(d_a);
  cudaFree(d_b);
  return 0;
}

我们通过定义主机和设备数组并为它们分配内存来启动了main函数。在第一个for循环中，h_a输入数据数组被初始化为从0到15的值。我们使用了取模运算，因此 1,000 个元素将均匀地分配到0到15的值之间。第二个数组，用于存储直方图，被初始化为零。这两个数组被复制到设备内存中。内核将计算直方图并将其返回到主机。我们将在控制台上打印这个直方图。输出如下所示：

当我们尝试使用原子操作来测量这段代码的性能并与 CPU 性能进行比较时，对于大型数组，它比 CPU 慢。这引发了一个问题：我们应该使用 CUDA 进行直方图计算，还是有可能使这种计算更快？

这个问题的答案是：是。如果我们为给定的块计算直方图使用共享内存，然后将这个块直方图添加到全局内存上的整体直方图中，那么可以加快操作速度。这是因为加法是一个累积操作。以下是用共享内存进行直方图计算的内核代码：

#include <stdio.h>
#include <cuda_runtime.h>
#define SIZE 1000
#define NUM_BIN 256
__global__ void histogram_shared_memory(int *d_b, int *d_a)
{
  int tid = threadIdx.x + blockDim.x * blockIdx.x;
  int offset = blockDim.x * gridDim.x;
  __shared__ int cache[256];
  cache[threadIdx.x] = 0;
  __syncthreads();

  while (tid < SIZE)
  {
    atomicAdd(&(cache[d_a[tid]]), 1);
    tid += offset;
  }
  __syncthreads();
  atomicAdd(&(d_b[threadIdx.x]), cache[threadIdx.x]);
}

在此代码中，桶的数量为 256 而不是 16，以提供更大的容量。我们定义的共享内存大小等于一个块中的线程数，即 256 个桶。我们将计算当前块的直方图，因此共享内存初始化为零，并按前面讨论的方式计算此块的直方图。但是，这次结果存储在共享内存中，而不是全局内存中。在这种情况下，只有 256 个线程试图访问共享内存中的 256 个内存元素，而不是像前一个代码中的 1,000 个元素。这将有助于减少原子操作中的时间开销。最后一行的最终原子add操作将一个块的直方图添加到整体直方图值中。由于加法是一个累积操作，我们不必担心每个块执行的顺序。此main函数与前面的函数类似。

这个内核函数的输出如下：

如果你测量前面程序的性能，它将击败没有共享内存的 GPU 版本和大型数组大小的 CPU 实现。你可以通过将 GPU 计算的直方图结果与 CPU 计算结果进行比较来检查 GPU 计算的直方图是否正确。

本节演示了在 GPU 上实现直方图的过程。它还强调了在 CUDA 程序中使用共享内存和原子操作的重要性。它还展示了 CUDA 在图像处理应用中的帮助以及将现有 CPU 代码转换为 CUDA 代码的简便性。

摘要

在本章中，我们看到了一些 CUDA 的高级概念，这些概念可以帮助我们使用 CUDA 开发复杂的应用程序。我们看到了测量设备代码性能的方法，以及如何使用 Nvidia Visual Profiler 工具查看内核函数的详细配置文件。这有助于我们识别降低程序性能的操作。我们看到了从 CUDA 代码本身处理硬件操作错误的方法，以及使用某些工具调试代码的方法。CPU 提供了有效的任务并行性，其中两个完全不同的函数可以并行执行。我们还看到 GPU 也通过 CUDA 流提供这种功能，并在相同的向量加法程序中使用 CUDA 流实现了两倍的速度提升。

然后，我们看到了使用 CUDA 加速排序算法的例子，这是构建复杂计算应用时需要理解的重要概念。图像处理是一个计算密集型任务，需要实时执行。几乎所有的图像处理算法都可以利用 GPU 和 CUDA 的并行性。因此，在最后一节中，我们看到了 CUDA 在加速图像处理应用中的应用，以及如何将现有的 C++代码转换为 CUDA 代码。我们还开发了用于直方图计算的 CUDA 代码，这是一个重要的图像处理应用。

本章也标志着与 CUDA 编程相关概念的结束。从下一章开始，我们将开始使用 OpenCV 库开发计算机视觉应用，该库利用了我们到目前为止所看到的 CUDA 加速概念。从下一章开始，我们将处理真实图像而不是矩阵。

问题

为什么不使用 CPU 计时器来衡量内核函数的性能？
尝试使用 Nvidia Visual Profiler 工具可视化上一章中实现的矩阵乘法代码的性能。
给出程序中遇到的不同语义错误示例。
内核函数中线程发散的缺点是什么？请用例子解释。
使用cudahostAlloc函数在主机上分配内存的缺点是什么？
证明以下陈述的正确性：CUDA 流中的操作顺序对于提高程序性能非常重要。
为了在 CUDA 中实现良好的性能，对于 1024 x 1024 的图像，应该启动多少个块和线程？

第五章：使用 CUDA 支持的 OpenCV 入门

到目前为止，我们已经看到了所有与使用 CUDA 进行并行编程相关的概念，以及它是如何利用 GPU 进行加速的。从本章开始，我们将尝试使用 CUDA 中的并行编程概念来应用于计算机视觉应用。虽然我们已经处理过矩阵，但我们还没有处理过实际图像。基本上，处理图像类似于二维矩阵的操作。我们不会从头开始为 CUDA 中的计算机视觉应用开发整个代码，但我们将使用名为 OpenCV 的流行计算机视觉库。尽管这本书假设读者对使用 OpenCV 有一些熟悉，但本章回顾了在 C++中使用 OpenCV 的概念。本章描述了在 Windows 和 Ubuntu 上安装支持 CUDA 的 OpenCV 库。然后它描述了如何测试此安装并运行一个简单的程序。本章通过为它开发简单的代码来描述使用 OpenCV 处理图像和视频。本章还将比较具有 CUDA 支持的程序与没有 CUDA 支持的程序的性能。

本章将涵盖以下主题：

图像处理和计算机视觉简介
支持 CUDA 的 OpenCV 简介
在 Windows 和 Ubuntu 上安装支持 CUDA 的 OpenCV
使用 OpenCV 处理图像
使用 OpenCV 处理视频
图像上的算术和逻辑运算
颜色空间转换和图像阈值
CPU 和 GPU OpenCV 程序的性能比较

技术要求

本章需要具备图像处理和计算机视觉的基本理解。它需要熟悉基本的 C 或 C++编程语言以及前几章中解释的所有代码示例。本章中使用的所有代码都可以从以下 GitHub 链接下载：github.com/PacktPublishing/Hands-On-GPU-Accelerated-Computer-Vision-with-OpenCV-and-CUDA。代码可以在任何操作系统上执行，尽管它只在 Ubuntu 16.04 上进行了测试。

查看以下视频以查看代码的实际运行情况：

bit.ly/2xF5cQV

图像处理和计算机视觉简介

世界上的图像和视频数据量每天都在不断增加。随着移动设备用于捕捉图像和互联网用于发布图像的日益普及，每天都会产生大量的视频和图像数据。图像处理和计算机视觉在各个领域的许多应用中都得到了应用。医生使用 MRI 和 X 射线图像进行医学诊断。空间科学家和化学工程师使用图像进行太空探索和分子水平上各种基因的分析。图像可以用于开发自动驾驶车辆和视频监控应用。它们还可以用于农业应用和制造过程中的故障产品识别。所有这些应用都需要在计算机上以高速处理图像。我们不会探讨图像是如何通过相机传感器捕捉并转换为计算机存储的数字图像的。在这本书中，我们只涵盖计算机上的图像处理，假设它已经存储好了。

许多人将图像处理和计算机视觉这两个术语互换使用。然而，这两个领域之间是有区别的。图像处理关注通过修改像素值来提高图像的视觉质量，而计算机视觉关注从图像中提取重要信息。因此，在图像处理中，输入和输出都是图像，而在计算机视觉中，输入是图像，但输出是从该图像中提取的信息。两者都有广泛的应用，但图像处理主要在计算机视觉应用的预处理阶段使用。

图像以多维矩阵的形式存储。因此，在计算机上处理图像不过是操作这个矩阵。我们在前面的章节中看到了如何在 CUDA 中处理矩阵。在 CUDA 中读取、操作和显示图像的代码可能会变得非常长、繁琐且难以调试。因此，我们将使用一个包含所有这些功能 API 的库，并且可以利用 CUDA-GPU 加速处理图像的优势。这个库被称为 OpenCV，它是“开放计算机视觉”的缩写。在下一节中，我们将详细介绍这个库。

OpenCV 简介

OpenCV 是一个以计算效率为前提，并专注于实时性能的计算机视觉库。它用 C/C++编写，包含超过一百个有助于计算机视觉应用的功能。OpenCV 的主要优势是它是开源的，并且根据伯克利软件发行许可（BSD 许可）发布，这允许在研究和商业应用中免费使用 OpenCV。这个库提供了 C、C++、Java 和 Python 语言的接口，并且可以在所有操作系统上使用，如 Windows、Linux、macOS 和 Android，而无需修改任何一行代码。

这个库还可以利用多核处理、OpenGL 和 CUDA 进行并行处理。由于 OpenCV 轻量级，它也可以在树莓派等嵌入式平台上使用。这使得它在现实场景中部署计算机视觉应用到嵌入式系统中变得理想。我们将在接下来的几章中探讨这一点。这些特性使 OpenCV 成为计算机视觉开发者的默认选择。它拥有广泛的开发者和用户社区，不断帮助改进这个库。OpenCV 的下载量以百万计，并且每天都在增加。另一个流行的计算机视觉和图像处理工具是 MATLAB，因此你可能会想知道使用 OpenCV 而不是 MATLAB 的优势。以下表格显示了这两个工具的比较：

参数	OpenCV	MATLAB
程序速度	由于它是用 C/C++开发的，所以速度更快	低于 OpenCV
资源需求	OpenCV 是一个轻量级库，因此在硬盘和 RAM 方面都消耗很少的内存。一个普通的 OpenCV 程序将需要少于 100MB 的 RAM。	MATLAB 非常庞大。最新版本的 MATLAB 安装可能需要在硬盘上占用超过 15GB 的空间，并且在使用时需要大量的 RAM（超过 1GB）。
可移植性	OpenCV 可以在所有可以运行 C 语言的操作系统上运行。	MATLAB 只能在 Windows、Linux 和 MAC 上运行。
成本	在商业或学术应用中使用 OpenCV 是完全免费的。	MATLAB 是许可软件，因此你必须支付一大笔钱才能在学术或商业应用中使用它。
易用性	与其较少的文档和难以记忆的语法相比，OpenCV 相对较难使用。它也没有自己的开发环境。	MATLAB 有自己的集成开发环境，内置帮助资源，这使得新程序员容易使用。

MATLAB 和 OpenCV 都有自己的优缺点。但当我们想在嵌入式应用中使用计算机视觉并利用并行处理时，OpenCV 是理想的选择。因此，在这本书中，我们将描述如何使用 GPU 和 CUDA 加速计算机视觉应用。OpenCV 提供了 C、C++、Python 和 Java 的 API。它是用 C/C++编写的，因此这些语言的 API 将是最快的。此外，CUDA 加速在 C/C++ API 中支持得更好，所以在这本书中，我们将使用 OpenCV 的 C/C++ API。在下一节中，我们将看到如何在各种操作系统上安装 OpenCV。

支持 CUDA 的 OpenCV 安装

使用 CUDA 安装 OpenCV 并不像你想象中那么简单。它涉及许多步骤。在本节中，我们将通过截图详细解释在 Windows 和 Ubuntu 上安装 OpenCV 的所有步骤，以便你可以轻松地设置你的环境。

在 Windows 上安装 OpenCV

本节解释了在 Windows 操作系统上安装带有 CUDA 的 OpenCV 所需的步骤。这些步骤在 Windows 10 操作系统上执行，但它们可以在任何 Windows 操作系统上工作。

使用预构建的二进制文件

对于 OpenCV，有一些预构建的二进制文件可供下载并直接用于您的程序中。它没有充分利用 CUDA，因此不建议在本书中使用。以下步骤描述了在 Windows 上不使用 CUDA 支持安装 OpenCV 的过程：

确保已安装 Microsoft Visual Studio 以编译 C 程序。
从sourceforge.net/projects/opencvlibrary/files/opencv-win/下载 OpenCV 的最新版本。
双击下载的.exe文件，并将其提取到您选择的文件夹中。在这里，我们将其提取到C://opencv文件夹。
通过右键单击我的电脑 | 高级设置 | 环境变量 | 新建来设置环境变量OPENCV_DIR。将其值设置为C:\opencv\build\x64\vc14，如下截图所示。这里vc14将取决于 Microsoft Visual Studio 的版本：

现在，您可以使用此安装为 OpenCV 应用程序使用 C/C++。

从源代码构建库

如果您想编译带有 CUDA 支持的 OpenCV，请按照以下步骤进行安装：

OpenCV 带有 CUDA 将需要一个 C 编译器和 GPU 编译器。它需要 Microsoft Visual Studio 和最新的 CUDA 安装。安装它们的步骤在第一章，介绍 CUDA 和 CUDA 入门中有所介绍。因此，在继续之前，请检查它们是否已正确安装。
通过访问链接下载 OpenCV 最新版本的源代码：github.com/opencv/opencv/.
有些额外的模块不包括在 OpenCV 中，但它们可以在名为opencv_contrib的额外模块中找到，这个模块可以与 OpenCV 一起安装。此模块中可用的功能不稳定；一旦它们变得稳定，它们就会被移动到实际的 OpenCV 源中。如果您想安装此模块，请从github.com/opencv/opencv_contrib下载。
从以下链接安装cmake：cmake.org/download/。它是编译 OpenCV 库所需的。
将opencv和opencv_contrib的 ZIP 文件提取到任何文件夹中。在这里，它们被提取到C://opencv和C://opencv_contrib文件夹中。
打开 CMake 以编译 OpenCV。在 CMake 中，您需要选择 OpenCV 源代码的路径，并选择此源代码将被构建的文件夹。如下截图所示：

然后点击配置。它将开始配置源。CMake 将尝试根据系统变量的路径设置定位尽可能多的包。配置过程如下面的截图所示：

如果某些包没有找到，您可以手动定位它们。为了配置具有 CUDA 支持的 OpenCV 安装，您必须检查如下截图所示的WITH_CUDA变量，然后再次点击配置：

配置完成后，点击生成。这将根据您选择的 Visual Studio 版本创建 Visual Studio 项目文件。当生成完成后，窗口应该类似于以下截图：

前往opencv文件夹的构建目录，并找到名为OpenCV.sln的 Visual Studio 项目，如下面的截图所示：

这将在 Microsoft Visual Studio 中打开项目。在解决方案资源管理器中，找到名为ALL_BUILD的项目。右键单击它并构建。在 Visual Studio 中为调试和发布选项构建此项目。如下面的截图所示：

构建整个项目可能需要很长时间，尽管这会根据您的处理器和 Visual Studio 版本而有所不同。在构建操作成功完成后，您就可以在 C/C++项目中使用 OpenCV 库了。
通过右键单击我的电脑 | 高级系统设置 | 环境变量 | 新建来设置环境变量OPENCV_DIR。将其值设置为C:\opencv\build\x64\vc14。在这里，vc14将取决于您使用的 Microsoft Visual Studio 版本：

您可以通过访问C://opencv/build/bin/Debug目录并运行任何.exe 应用程序来检查安装情况。

在 Linux 上安装具有 CUDA 支持的 OpenCV

本节涵盖了在 Linux 操作系统上安装具有 CUDA 支持的 OpenCV 的步骤。这些步骤已在 Ubuntu 16.04 上测试，但它们应该适用于任何 Unix 发行版：

具有 CUDA 的 OpenCV 将需要最新的 CUDA 安装。安装它的过程在第一章，介绍 CUDA 和 CUDA 入门中有所覆盖。所以在继续之前，请检查它是否已正确安装。您可以通过执行nvidia-smi命令来检查 CUDA 工具包和 Nvidia 设备驱动程序的安装情况。如果您的安装正常工作，您应该看到如下类似的输出：

通过访问链接下载 OpenCV 最新版本的源代码：github.com/opencv/opencv/。将其解压到opencv文件夹中。
有一些额外的模块不包括在 OpenCV 中，但它们在名为opencv_contrib的额外模块中可用，可以与 OpenCV 一起安装。此模块中提供的函数是不稳定的；一旦它们变得稳定，它们就会被移动到实际的 OpenCV 源中。如果您想安装此模块，可以从github.com/opencv/opencv_contrib下载它。将其提取到与opencv文件夹相同的目录中的opencv_contrib文件夹。
打开opencv文件夹并创建一个构建目录。然后进入这个新创建的build目录。这些步骤可以通过在命令提示符中执行以下命令来完成：

$ cd opencv
$ mkdir build
$ cd build

cmake命令用于编译具有 CUDA 支持的opencv。确保在此命令中将WITH_CUDA标志设置为 ON，并指定一个适当的路径，用于存储在opencv_contrib目录中下载并保存的额外模块。完整的cmake命令如下所示：

cmake -D CMAKE_BUILD_TYPE=RELEASE CMAKE_INSTALL_PREFIX=/usr/local WITH_CUDA=ON  ENABLE_FAST_MATH=1 CUDA_FAST_MATH=1 -D WITH_CUBLAS=1 OPENCV_EXTRA_MODULES_PATH=../../opencv_contrib/modules BUILD_EXAMPLES=ON ..

它将开始配置和创建makefile。它将根据系统路径中的值定位所有额外模块。以下截图显示了具有所选 CUDA 安装的cmake命令的输出：

在配置成功后，CMake 将在构建目录中创建一个 makefile。要使用此 makefile 编译 OpenCV，请在命令窗口中执行命令make -j8，如下所示：

编译成功后，要安装 OpenCV，您必须从命令行执行命令sudo make install。以下是该命令的输出：

运行sudo ldconfig命令以完成安装。它创建了必要的链接和缓存到opencv库。
您可以通过运行opencv/samples/gpu文件夹中的任何示例来检查安装。

要在程序中使用 OpenCV，您必须包含opencv2/opencv.hpp头文件。此头文件将包含程序所需的全部其他头文件。因此，所有 OpenCV 程序都必须在顶部包含此头文件。

在 OpenCV 中处理图像

现在 OpenCV 已安装在系统上，我们可以开始使用它来处理图像。在本节中，我们将学习 OpenCV 中图像的表示方式，开发读取图像、显示图像和将图像保存到磁盘的程序。我们还将了解在 OpenCV 中创建合成图像的方法。我们还将使用 OpenCV 在图像上绘制不同的形状。此外，还将解释 OpenCV 的重要语法和功能。

OpenCV 中的图像表示

如前所述，图像不过是二维数组，因此它们应该在计算机内部以数组的形式存储以供处理。OpenCV 提供了一个 Mat 类，它实际上是一个用于存储图像的图像容器。Mat 对象可以通过以下两行分别创建并分配给图像：

Mat img;
img= imread("cameraman.tif");

在创建对象时，也可以定义图像的数据类型和二维数组的大小。图像的数据类型非常重要，因为它表示了指定单个像素值所使用的通道数和位数。灰度图像只有一个通道，而彩色图像是三个独立通道的组合：红色、绿色和蓝色。

单个像素使用的位数指定了离散灰度级别值的数量。一个 8 位图像可以具有介于 0 到 255 之间的灰度级别，而 16 位图像可以具有介于 0 到 65,535 之间的灰度级别。OpenCV 支持许多数据类型，默认为 CV_8U，表示单通道的 8 位无符号图像。它等同于 CV_8UC1。彩色图像可以指定为 CV_8UC3，表示具有三个通道的 8 位无符号图像。OpenCV 支持多达 512 个通道。五个或更多通道必须用圆括号定义，例如，CV_8UC(5) 表示具有五个通道的 8 位图像。OpenCV 还支持有符号数，因此数据类型也可以是 CV_16SC3，表示具有三个通道的 16 位有符号图像。

Mat 对象可以用来定义图像的大小。这也被称为图像的分辨率。它表示水平和垂直方向上的像素数量。通常，图像的分辨率是以宽度 x 高度来定义的。而在 Mat 对象中的数组大小应该以行 x 列数来定义。以下是一些使用 Mat 定义图像容器的示例：

Mat img1(6,6,CV_8UC1); 
//This defines img1 object with size of 6x6, unsigned 8-bit integers and single channel.

Mat img2(256,256, CV_32FC1)
//This defines img2 object with size of 256x256, 32 bit floating point numbers and single channel.

Mat img3(1960,1024, CV_64FC3)
//This defines img3 object with size of 1960x1024, 64 bit floating point numbers and three channels.

图像的分辨率和大小将决定在磁盘上保存图像所需的空间。假设一个具有三个通道的彩色图像的大小为 1024 x 1024，那么它将占用 3 x 1024 x 1024 字节 = 3 MB 的磁盘空间。在下一节中，我们将看到如何使用这个 Mat 对象和 OpenCV 读取和显示图像。

读取和显示图像

在本节中，我们将尝试使用 C++ 和 OpenCV 开发读取和显示图像的代码。整个代码如下，然后逐行进行解释：

#include <opencv2/opencv.hpp>
#include <iostream>

using namespace cv;
using namespace std;

int main(int argc, char** argv)
{
  // Read the image 
  Mat img = imread("images/cameraman.tif",0);

  // Check for failure in reading an Image
  if (img.empty()) 
  {
    cout << "Could not open an image" << endl;
    return -1;
  }
  //Name of the window
  String win_name = "My First Opencv Program"; 

  // Create a window
  namedWindow(win_name); 

  // Show our image inside the created window.
  imshow(win_name, img); 

  // Wait for any keystroke in the window 
  waitKey(0); 

  //destroy the created window
  destroyWindow(win_name); 

  return 0;
}

程序从包含标准输入输出和图像处理的头文件开始。

程序中使用来自 std 命名空间的功能，如 cout 和 endl，因此添加了 std 命名空间。所有 OpenCV 类和函数都是使用 cv 命名空间定义的。因此，为了使用 cv 命名空间中定义的函数，我们指定了 using namespace cv 这一行。如果省略该行，那么在 cv 命名空间中的每个函数都必须以下述方式使用：

Mat img = cv::imread("cameraman.tif")

main函数包含读取和显示图像的代码。在 OpenCV 中使用imread命令读取图像。它返回一个Mat对象。imread命令有两个参数。第一个参数是图像的名称及其路径。路径可以有两种指定方式。您可以在 PC 上指定图像的完整路径，或者从您的代码文件中指定图像的相对路径。在上面的示例中，使用了相对路径，其中图像位于与代码文件相同的目录下的 images 文件夹中。

第二个参数是可选的，用于指定图像是要以灰度图像还是彩色图像读取。如果图像要以彩色图像读取，则指定IMREAD_COLOR或1。如果图像要以灰度图像读取，则指定IMREAD_GRAYSCALE或0。如果要以保存的格式读取图像，则将IMREAD_UNCHANGED或-1作为第二个参数。如果图像作为彩色图像读取，imread命令将返回三个通道，从蓝色、绿色和红色（BGR格式）开始。如果第二个参数未提供，则默认值为IMREAD_COLOR，它将图像作为彩色图像读取。

如果图像无法读取或磁盘上不可用，则imread命令将返回一个Null Mat对象。如果发生这种情况，无需继续执行进一步图像处理代码，我们可以在此时通过通知用户关于错误来退出。这由if循环内的代码处理。

应创建一个用于显示图像的窗口。OpenCV 提供了一个名为namedWindow的函数来实现这一功能。它需要两个参数。第一个参数是窗口的名称，它必须是一个字符串。第二个参数指定要创建的窗口的大小。它可以取两个值：WINDOW_AUTOSIZE或WINDOW_NORMAL。如果指定了WINDOW_AUTOSIZE，则用户将无法调整窗口大小，图像将以原始大小显示。如果指定了WINDOW_NORMAL，则用户可以调整窗口大小。此参数是可选的，如果未指定，其默认值为WINDOW_AUTOSIZE。

要在创建的窗口中显示图像，使用imshow命令。此命令需要两个参数。第一个参数是使用namedWindow命令创建的窗口名称，第二个参数是要显示的图像变量。此变量必须是一个Mat对象。要显示多个图像，必须创建具有唯一名称的单独窗口。窗口的名称将作为图像窗口的标题出现。

imshow 函数应该有足够的时间在创建的窗口中显示一张图片。这是通过使用 waitKey 函数来实现的。因此，在所有 OpenCV 程序中，imshow 函数后面都应该跟 waitkey 函数，否则图片将不会显示。waitKey 是一个键盘绑定函数，它接受一个参数，即以毫秒为单位的时间。它将在指定的时间内等待按键，然后移动到下一行代码。如果没有指定参数或指定为 0，它将无限期地等待按键。只有在键盘上按下任何键时，它才会移动到下一行。我们还可以检测是否按下了特定的键，并根据按下的键做出某些决定。我们将在本章后面使用这个功能。

在程序终止之前，需要关闭为显示窗口创建的所有窗口。这可以通过使用 destroyAllWindows 函数来完成。它将关闭程序中通过 namedWindow 函数创建的所有窗口。有一个名为 destroyWindow 的函数，可以关闭特定的窗口。应将窗口名称作为参数提供给 destroyWindow 函数。

对于程序的执行，只需将代码复制并粘贴到 Windows 上的 Visual Studio 中，或者如果你在 Ubuntu 上使用，可以创建一个 cpp 文件。构建方法与 Visual Studio 中的正常 cpp 应用程序类似，所以这里不再重复。要在 Ubuntu 上执行，请在保存 cpp 文件的文件夹中从命令提示符执行以下命令：

For compilation:
$ g++ -std = c++11 image_read.cpp 'pkg_config --libs --cflags opencv' -o image_read 
For execution:
$./image_read

前面程序的输出如下：

读取和显示彩色图像

在前面的程序中，imread 的第二个参数被指定为 0，这意味着它将以灰度图像的形式读取图像。假设你想读取任何彩色图像。为此，你可以按以下方式更改 imread 命令：

Mat img = imread("images/autumn.tif",1);

第二个参数被指定为 1，这意味着它将以 BGR 格式读取图像。重要的是要注意，OpenCV 的 imread 和 imshow 使用 BGR 格式来处理彩色图像，这与 MATLAB 和其他图像处理工具使用的 RGB 格式不同。更改 imread 后的输出如下：

即使是彩色图像，也可以通过将第二个参数指定为 0 来将其读取为灰度图像。这将隐式地将图像转换为灰度并读取。图像将看起来如下：

记住如何使用 imread 函数读取图像非常重要，因为它将影响你程序中的其他图像处理代码。

总结来说，在本节中，我们看到了如何使用 OpenCV 读取图像并显示它。在这个过程中，我们还了解了一些 OpenCV 中可用的重要函数。在下一节中，我们将看到如何使用 OpenCV 创建合成图像。

使用 OpenCV 创建图像

有时，我们可能需要创建自己的图像或在现有图像上绘制一些形状。或者，我们可能想在检测到的对象周围绘制边界框或在图像上显示标签。因此，在本节中，我们将看到如何创建空白灰度和彩色图像。我们还将看到在图像上绘制线条、矩形、椭圆、圆和文本的函数。

要创建一个大小为 256 x 256 的空黑色图像，可以使用以下代码：

#include <opencv2/opencv.hpp>
#include <iostream>

using namespace cv;
using namespace std;

int main(int argc, char** argv)
{
  //Create blank black grayscale Image with size 256x256
  Mat img(256, 256, CV_8UC1, Scalar(0)); 
  String win_name = "Blank Image"; 
  namedWindow(win_name); 
  imshow(win_name, img); 
  waitKey(0); 
  destroyWindow(win_name); 
  return 0;
}

代码与用于读取图像的代码大致相似，但在这里不是使用 imread 命令，而是仅使用 Mat 类的构造函数来创建图像。如前所述，我们可以在创建 Mat 对象时提供大小和数据类型。因此，在创建 img 对象时，我们提供了四个参数。前两个参数指定了图像的大小，第一个参数定义了行数（高度），第二个参数定义了列数（宽度）。第三个参数定义了图像的数据类型。我们使用了 CV_8UC1，这意味着一个单通道的 8 位无符号整数图像。最后一个参数指定了数组中所有像素的初始化值。

这里我们使用了 0，这是黑色的值。当程序执行时，它将创建一个大小为 256 x 256 的黑色图像，如下所示：

可以使用类似的代码创建任何颜色的空白图像，如下所示：

#include <opencv2/opencv.hpp>
#include <iostream>

using namespace cv;
using namespace std;

int main(int argc, char** argv)
{
  //Create blank blue color Image with size 256x256
  Mat img(256, 256, CV_8UC3, Scalar(255,0,0)); 
  String win_name = "Blank Blue Color Image"; 
  namedWindow(win_name); 
  imshow(win_name, img); 
  waitKey(0); 
  destroyWindow(win_name); 
  return 0;
}

在创建 Mat 对象时，不是使用 CV_8UC1 数据类型，而是使用 CV_8UC3，它指定了一个具有三个通道的 8 位图像。因此，单个像素有 24 位。第四个参数指定了起始像素值。它使用标量关键字和一个包含所有三个通道起始值的元组来指定。在这里，蓝色通道初始化为 255，绿色通道初始化为 0，红色通道初始化为 0。这将创建一个大小为 256 x 256 的蓝色图像。元组中值的组合将创建不同的颜色。前述程序的输出如下：

在空白图像上绘制形状

要开始在图像上绘制不同的形状，我们将首先使用以下命令创建一个任意大小的空白黑色图像：

Mat img(512, 512, CV_8UC3, Scalar(0,0,0));

此命令将创建一个大小为 512 x 512 的黑色图像。现在，我们将开始在这个图像上绘制不同的形状。

绘制线条

一条线可以通过两个点来指定：起点和终点。要在图像上绘制线条，必须指定这两个点。在图像上绘制线条的函数如下：

line(img,Point(0,0),Point(511,511),Scalar(0,255,0),7);

线函数有五个参数。第一个参数指定需要绘制线的图像，第二个和第三个参数分别定义起点和终点。这些点使用 Point 类构造函数定义，该构造函数以图像的 x 和 y 坐标为参数。第四个参数指定线的颜色。它指定为 B、G 和 R 值的元组。这里，取值为 (0,255,0)，指定绿色。第五个参数是线的厚度。其值被设置为 7 像素宽。此函数还有一个可选的 linetype 参数。前面的函数将绘制一个从 (0,0) 到 (511,511) 的对角线绿色线，宽度为 7 像素。

绘制矩形

矩形可以使用两个极端对角点来指定。OpenCV 提供了一个在图像上绘制矩形的函数，其语法如下：

rectangle(img,Point(384,0),Point(510,128),Scalar(255,255,0),5);

矩形函数有五个参数。第一个参数是需要绘制矩形的图像。第二个参数是矩形的左上角点。第三个参数是矩形的右下角点。第四个参数指定边框的颜色。它指定为 (255,255,0)，是蓝色和绿色的混合，产生青色。第五个参数是边框的厚度。如果第五个参数被指定为 -1，则形状将被填充。因此，前面的函数将使用两个极端点 (384,0) 和 (510,128)，以青色绘制一个边框厚度为 5 像素的矩形。

绘制圆

圆可以通过中心和半径来指定。OpenCV 提供了一个在图像上绘制圆的函数，其语法如下：

circle(img,Point(447,63), 63, Scalar(0,0,255), -1);

圆函数有五个参数。第一个参数是需要绘制圆的图像。第二个参数指定该圆的中心点，第三个参数指定半径。第四个参数指定圆的颜色。取值为 (0,0,255)，表示红色。第五个参数是边框的厚度。这里，它被设置为 -1，意味着圆将被红色填充。

绘制椭圆

OpenCV 提供了一个在图像上绘制椭圆的函数，其语法如下：

ellipse(img,Point(256,256),Point(100,100),0,0,180,255,-1);

椭圆函数有许多参数。第一个参数指定需要绘制椭圆的图像。第二个参数指定椭圆的中心。第三个参数指定椭圆将要绘制的框的大小。第四个参数指定椭圆需要旋转的角度。它被设置为 0 度。第五个和第六个参数指定椭圆需要绘制的角度范围。它被设置为 0 到 180 度。因此，只会绘制椭圆的一半。下一个参数指定椭圆的颜色，它被指定为 255. 它与 (255,0,0) 相同，表示蓝色。最后一个参数指定边框的粗细。它被设置为 -1，因此椭圆将被蓝色填充。

在图像上写文本

OpenCV 提供了一个在图像上写文本的函数，即 putText。该函数的语法如下：

putText( img, "OpenCV!", Point(10,500), FONT_HERSHEY_SIMPLEX, 3,Scalar(255, 255, 255), 5, 8 );

putText 函数有许多参数。第一个参数是要在图像上写文本的图像。第二个参数是作为字符串数据类型的文本，我们想要在图像上写的内容。第三个参数指定文本的左下角。第四个参数指定字体类型。OpenCV 中有许多可用的字体类型，你可以查看 OpenCV 文档。第五个参数指定字体的缩放比例。第六个参数是文本的颜色。它被设置为 (255,255,255)，表示白色。第七个参数是文本的粗细，它被设置为 5，最后一个参数指定线型，它被设置为 8。

我们已经看到了在空黑图像上绘制形状的单独函数。以下代码显示了之前讨论的所有函数的组合：

#include <opencv2/opencv.hpp>
#include <iostream>

using namespace cv;
using namespace std;

int main(int argc, char** argv)
{

  Mat img(512, 512, CV_8UC3, Scalar(0,0,0)); 
  line(img,Point(0,0),Point(511,511),Scalar(0,255,0),7);
  rectangle(img,Point(384,0),Point(510,128),Scalar(255,255,0),5);
  circle(img,Point(447,63), 63, Scalar(0,0,255), -1);
  ellipse(img,Point(256,256),Point(100,100),0,0,180,255,-1);
  putText( img, "OpenCV!", Point(10,500), FONT_HERSHEY_SIMPLEX, 3,Scalar(255, 255,  255), 5, 8 );
  String win_name = "Shapes on blank Image"; 
  namedWindow(win_name); 
  imshow(win_name, img); 
  waitKey(0); 
  destroyWindow(win_name); 
  return 0;
}

上一段代码的输出图像如下：

将图像保存到文件

图像也可以从 OpenCV 程序保存到磁盘。当我们想要将处理后的图像存储到计算机上的磁盘时，这很有用。OpenCV 提供了 imwrite 函数来完成此操作。此函数的语法如下：

bool flag = imwrite("images/save_image.jpg", img);

imwrite 函数接受两个参数。第一个参数是你想要保存的文件名及其路径。第二个参数是你想要保存的 img 变量。此函数返回一个布尔值，表示文件是否成功保存到磁盘上。

在本节中，我们使用 OpenCV 处理了图像。在下一节中，我们将使用 OpenCV 处理视频，视频不过是图像的序列。

在 OpenCV 中处理视频

本节将展示使用 OpenCV 从文件和摄像头读取视频的过程。它还将描述将视频保存到文件的过程。这也可以与连接到计算机的 USB 摄像头一起工作。视频不过是图像的序列。尽管 OpenCV 不是针对视频处理应用程序优化的，但它在这方面做得相当不错。OpenCV 无法捕获音频，因此我们必须使用其他一些与 OpenCV 一起使用的工具来捕获音频和视频。

在计算机上处理视频

本节描述了读取存储在计算机上的视频文件的过程。在所有使用 OpenCV 的视频处理应用程序中，视频的所有帧将逐个读取、处理并在屏幕上显示。

以下代码用于读取和显示视频——随后将逐行解释：

#include <opencv2/opencv.hpp>
#include <iostream>
using namespace cv;
using namespace std;
int main(int argc, char* argv[])
{
  //open the video file from PC
  VideoCapture cap("images/rhinos.avi"); 
  // if not success, exit program
  if (cap.isOpened() == false) 
  {
    cout << "Cannot open the video file" << endl;
    return -1;
  }
  cout<<"Press Q to Quit" << endl;
  String win_name = "First Video";
  namedWindow(win_name); 
  while (true)
  {
    Mat frame;
    // read a frame
    bool flag = cap.read(frame); 

    //Breaking the while loop at the end of the video
    if (flag == false) 
    {
      break;
    }
    //display the frame 
    imshow(win_name, frame);
    //Wait for 100 ms and key 'q' for exit
    if (waitKey(100) == 'q')
    {
      break;
    }
  }
  destroyWindow(win_name);
  return 0;
}

在包含库之后，在主函数中处理视频的第一件事是创建一个VideoCapture对象。VideoCapture类提供了许多构造函数，可用于处理视频。当我们想要处理存储在计算机上的视频文件时，我们需要在创建VideoCapture对象时将视频的名称及其路径作为参数传递给构造函数。

此对象提供了许多方法和属性，可以提供与视频相关的信息。我们将根据需要查看它们。它提供了isopened属性，该属性指示对象创建是否成功以及视频是否可用。它返回一个布尔值。如果cap.isopened为false，则视频不可用，因此无需在程序中进一步操作。这通过一个if循环处理，当视频不可用时，会通知用户并退出程序。

VideoCapture类提供了一个读取方法，可以逐个捕获帧。为了处理整个视频，我们必须启动一个连续的循环，该循环一直运行到视频的末尾。无限while循环可以完成这项工作。在while循环内部，使用读取方法读取第一帧。此方法有一个参数。它是一个 Mat 对象，我们希望在其中存储帧。它返回一个布尔值，指示帧是否已成功读取。当循环到达视频的末尾时，此布尔值将返回false，表示没有可用的帧。这个标志在循环中持续检查视频的结束；如果检测到，我们将使用break语句退出while循环。

帧是单个图像，因此显示该帧的过程与之前看到的是相同的。在上面的代码中，waitKey函数在if语句中使用。它会在每个帧之后等待 100 毫秒以获取按键。if语句正在检查按键是否为q。如果是q，则表示用户想要退出视频，因此在if中包含了 break 语句。

此代码将在整个视频播放完毕或用户在键盘上按下q键时终止视频显示。在整个本书中，我们将使用这种编码实践来处理视频。前一个程序的输出如下。截图是视频的一帧：

我们在每帧之间使用了 100 毫秒的延迟。当你将这个值减少到，比如，10 毫秒时，你认为会发生什么？答案是，每帧将显示得更快。这并不意味着视频的帧率改变了。这只是意味着帧之间的延迟减少了。如果你想看到视频的实际帧率，可以使用cap对象的CAP_PROP_FPS属性。它可以使用以下代码显示：

double frames_per_second = cap.get(CAP_PROP_FPS); 
cout << "Frames per seconds of the video is : " << frames_per_second ;

cap对象还具有其他属性，例如CAP_PROP_FRAME_WIDTH和CAP_PROP_FRAME_HEIGHT，它们表示帧的宽度和高度。这些属性也可以通过get方法获取。这些属性可以通过使用 cap 对象的set方法来设置。set方法有两个参数。第一个参数是属性的name，第二个参数是我们想要设置的value。

本节描述了从文件中读取视频的方法。下一节将展示从网络摄像头或 USB 摄像头处理视频的过程。

处理网络摄像头视频

本节描述了从连接到计算机的网络摄像头或 USB 摄像头捕获视频的过程。OpenCV 的好处是，相同的代码将适用于笔记本电脑和任何可以运行 C/C++的嵌入式系统。这有助于在任意硬件平台上部署计算机视觉应用。捕获视频并显示的代码如下：

#include <opencv2/opencv.hpp>
#include <iostream>

using namespace cv;
using namespace std;

int main(int argc, char* argv[])
{
  //open the Webcam
  VideoCapture cap(0); 
  // if not success, exit program
  if (cap.isOpened() == false) 
  {
    cout << "Cannot open Webcam" << endl;
    return -1;
  }
  //get the frames rate of the video from webcam
  double frames_per_second = cap.get(CAP_PROP_FPS); 
  cout << "Frames per seconds : " << frames_per_second << endl;
  cout<<"Press Q to Quit" <<endl;
  String win_name = "Webcam Video";
  namedWindow(win_name); //create a window
  while (true)
  {
    Mat frame;
    bool flag = cap.read(frame); // read a new frame from video 
    //show the frame in the created window
    imshow(win_name, frame);
    if (waitKey(1) == 'q')
    {
      break;
    }
  }
  return 0;
}

在从网络摄像头或 USB 摄像头捕获视频时，需要将那个摄像头的设备 ID 作为VideoCapture对象构造函数的参数提供。连接的主摄像头将具有设备 ID 零。笔记本电脑的网络摄像头或 USB 摄像头（当没有网络摄像头时）将具有设备 ID 零。如果有多个摄像头连接到设备，它们的设备 ID 将是(0,1)，依此类推。在前面的代码中，零表示代码将使用主摄像头来捕获视频。

另一段代码与从文件读取视频的代码大致相似。在这里，视频的帧率也被获取并显示。帧将以每 1 毫秒的间隔逐个读取，并在创建的窗口中显示。您必须按q键来终止操作。使用网络摄像头捕获的视频输出如下：

将视频保存到磁盘

要从 OpenCV 程序保存视频，我们需要创建VideoWriter类的对象。将视频保存到文件的代码如下：

Size frame_size(640, 640);
int frames_per_second = 30;

VideoWriter v_writer("images/video.avi", VideoWriter::fourcc('M', 'J', 'P', 'G'), frames_per_second, frame_size, true); 

//Inside while loop
v_writer.write(frame); 

//After finishing video write
v_writer.release();

在创建VideoWriter类的对象时，构造函数接受五个参数。第一个参数是你想要保存的视频文件名及其绝对或相对路径。第二个参数是用于视频编解码器的四个字符代码。它是通过VideoWriter::fourcc函数创建的。在这里，我们使用运动 JPEG 编解码器，因此它的四个字符代码是'M'、'J'、'P'和'G'。根据你的需求和操作系统，可以使用其他编解码器。第三个参数是每秒帧数。它可以指定为一个之前定义的整数变量或直接在函数中的整数值。在前面的代码中，使用了每秒30帧。第四个参数是帧的大小。它使用size关键字和两个参数frame_width和frame_height定义。在前面的代码中，它被定义为 640 x 640。第五个参数指定要存储的帧是彩色还是灰度。如果是真的，帧将以彩色帧保存。

要开始使用VideoWriter对象写入帧，OpenCV 提供了一个write方法。此方法用于逐个将帧写入视频，因此它包含在一个无限while循环中。此方法仅接受一个参数，即帧变量的名称。帧的大小应与创建VideoWriter对象时指定的相同。在写入完成后，重要的是刷新并关闭创建的视频文件。这可以通过使用release方法释放创建的VideoWriter对象来完成。

总结来说，在本节中，我们探讨了从设备上的文件或摄像头读取视频的过程。我们还看到了将视频写入文件的代码。从下一节开始，我们将看到如何使用 CUDA 加速的 OpenCV 在图像或视频上操作。

使用 OpenCV CUDA 模块的基本计算机视觉应用

在前面的章节中，我们了解到 CUDA 提供了一个优秀的接口，可以充分利用 GPU 的并行计算能力来加速复杂的计算应用。在本节中，我们将看到如何利用 CUDA 的能力与 OpenCV 一起用于计算机视觉应用。

OpenCV CUDA 模块简介

OpenCV 有一个 CUDA 模块，它包含数百个可以利用 GPU 能力的函数。它仅支持 Nvidia GPU，因为它在后台使用 Nvidia CUDA 运行时。为了使用 CUDA 模块，OpenCV 必须编译时设置WITH_CUDA标志为 ON。

使用 OpenCV CUDA 模块的亮点之一是它提供了一个与常规 OpenCV API 相似的 API。它也不需要详细了解 CUDA 编程，尽管了解 CUDA 和 GPU 架构不会有害。研究人员已经表明，使用具有 CUDA 加速的函数可以比类似的 CPU 函数提供 5x-100x 的速度提升。

在下一节中，我们将看到如何使用 CUDA 模块与 OpenCV 结合，在各种计算机视觉和图像处理应用中使用，这些应用在图像的各个像素上操作。

图像上的算术和逻辑运算

在本节中，我们将看到如何在图像上执行各种算术和逻辑运算。我们将使用 OpenCV CUDA 模块中定义的函数来执行这些操作。

两个图像的相加

当两个图像大小相同时，可以执行两个图像的相加。OpenCV 在cv::cuda命名空间内提供了一个add函数用于加法操作。它执行两个图像的逐像素相加。假设在两个图像中，(0,0)处的像素强度值分别为 100 和 150。结果图像中的强度值将是 250，这是两个强度值的相加。OpenCV 的加法是一个饱和操作，这意味着如果加法的结果超过 255，它将被饱和到 255。执行加法的代码如下：

#include <iostream>
#include "opencv2/opencv.hpp"

int main (int argc, char* argv[])
{
  //Read Two Images 
  cv::Mat h_img1 = cv::imread("images/cameraman.tif");
  cv::Mat h_img2 = cv::imread("images/circles.png");
  //Create Memory for storing Images on device
  cv::cuda::GpuMat d_result1,d_img1, d_img2;
  cv::Mat h_result1;
  //Upload Images to device     
  d_img1.upload(h_img1);
  d_img2.upload(h_img2);

  cv::cuda::add(d_img1,d_img2, d_result1);
  //Download Result back to host
  d_result1.download(h_result1);
  cv::imshow("Image1 ", h_img1);
  cv::imshow("Image2 ", h_img2);
  cv::imshow("Result addition ", h_result1);
  cv::imwrite("images/result_add.png", h_result1);
  cv::waitKey();
  return 0;
}

当任何计算机视觉操作需要在 GPU 上执行时，图像必须存储在设备内存中。为此内存的分配可以使用gpumat关键字，它与用于主机内存的 Mat 类型相似。图像的读取方式与之前相同。读取两个图像用于相加，并存储在主机内存中。这些图像使用设备memory变量的upload方法复制到设备内存。主机图像变量作为参数传递给此方法。

GPU CUDA 模块中的函数定义在cv::cuda命名空间中。它需要设备内存中的图像作为其参数。CUDA 模块中的 add 函数用于图像相加。它需要三个参数。前两个参数是要相加的两个图像，最后一个参数是结果将存储的目标。所有三个变量都应该使用gpumat定义。

使用设备变量的download方法将结果图像复制回主机。将结果复制的img主机变量作为参数传递给download方法。然后使用上一节中解释的相同函数显示和存储此图像。程序的输出如下：

从两个图像中减去

使用 OpenCV 和 CUDA 可以在图像上执行其他算术运算。OpenCV 提供了subtract函数来减去两个图像。它也是一个饱和操作，这意味着当减法的结果低于零时，它将被饱和到零。subtract命令的语法如下：

//d_result1 = d_img1 - d_img2
cv::cuda::subtract(d_img1, d_img2,d_result1);

再次，两个要减去的图像作为前两个参数提供，结果图像作为第三个参数提供。两个图像之间的减法结果如下：

图像混合

有时需要以不同的比例混合两个图像，而不是直接将两个图像相加。图像混合可以用以下方程式表示：

result = α * img1 + β * img2 + γ

这可以通过 OpenCV 中的addWeighted函数轻松实现。该函数的语法如下：

cv::cuda::addWeighted(d_img1,0.7,d_img2,0.3,0,d_result1)

该函数有六个参数。第一个参数是第一个源图像，第二个参数是第一个图像的混合权重，第三个参数是第二个源图像，第四个参数是第二个图像的混合权重，第五个参数是在混合时需要添加的常数伽玛，最后一个参数指定结果需要存储的目标位置。该函数将img1的 70%和img2的 30%用于混合。该函数的输出如下：

图像反转

除了算术运算外，OpenCV 还提供了对单个位进行操作的布尔运算。它包括 AND、OR、NOT 等。AND 和 OR 在掩码操作中非常有用，我们将在后面看到。NOT 操作用于反转图像，其中黑色转换为白色，白色转换为黑色。它可以表示为以下方程式：

result_image = 255 - input_image

在方程式中，255表示 8 位图像的最大强度值。进行图像反转的程序如下：

#include <iostream>
#include "opencv2/opencv.hpp"

int main (int argc, char* argv[])
{
  cv::Mat h_img1 = cv::imread("images/circles.png");
  //Create Device variables
  cv::cuda::GpuMat d_result1,d_img1;
  cv::Mat h_result1;     
  //Upload Image to device
  d_img1.upload(h_img1);

  cv::cuda::bitwise_not(d_img1,d_result1);

  //Download result back  to host
  d_result1.download(h_result1);
  cv::imshow("Result inversion ", h_result1);
  cv::imwrite("images/result_inversion.png", h_result1);
  cv::waitKey();
  return 0;
}

该程序与算术运算程序类似。使用bitwise_not函数进行图像反转。图像应该是灰度图像。它接受两个参数。第一个参数指示要反转的源图像，第二个参数指示反转图像要存储的目标位置。bitwise_not操作的输出如下：

如所示，通过反转，白色转换为黑色，黑色转换为白色。

总结来说，在本节中，我们看到了使用 OpenCV 和 CUDA 的各种算术和逻辑操作。在下一节中，我们将看到一些在计算机视觉应用中广泛使用的计算机视觉操作。

改变图像的色彩空间

如前所述，OpenCV 可以以灰度图像或彩色图像的形式读取图像，彩色图像具有三个通道：绿色、蓝色和红色，这种格式称为 BGR 格式。其他图像处理软件和算法在 RGB 图像上工作，其中红色通道后面是绿色和蓝色。还有许多其他颜色格式可以用于特定应用。这些包括 HSV 颜色空间，其中三个通道是色调、饱和度和亮度。色调代表颜色值，饱和度表示颜色的灰度级别，亮度代表颜色的亮度。另一个颜色空间是 YCrCb，它也非常有用。这个系统用图像中的一个亮度分量：亮度(Y)，和两个色度分量：色度(Cb 和 Cr)来表示颜色。

OpenCV 支持许多其他颜色空间，例如 XYZ、HLS、Lab 等。OpenCV 支持超过 150 种颜色转换方法。使用 OpenCV 中的cvtColor函数可以将一种颜色空间转换为另一种颜色空间。以下是一个使用此函数在不同颜色空间之间转换的示例：

#include <iostream>
#include "opencv2/opencv.hpp"

int main (int argc, char* argv[])
{
  cv::Mat h_img1 = cv::imread("images/autumn.tif");
  //Define device variables
  cv::cuda::GpuMat d_result1,d_result2,d_result3,d_result4,d_img1;
  //Upload Image to device
  d_img1.upload(h_img1);

  //Convert image to different color spaces
  cv::cuda::cvtColor(d_img1, d_result1,cv::COLOR_BGR2GRAY);
  cv::cuda::cvtColor(d_img1, d_result2,cv::COLOR_BGR2RGB);
  cv::cuda::cvtColor(d_img1, d_result3,cv::COLOR_BGR2HSV);
  cv::cuda::cvtColor(d_img1, d_result4,cv::COLOR_BGR2YCrCb);

  cv::Mat h_result1,h_result2,h_result3,h_result4;
  //Download results back to host
  d_result1.download(h_result1);
  d_result2.download(h_result2);
  d_result3.download(h_result3);
  d_result4.download(h_result4);

  cv::imshow("Result in Gray ", h_result1);
  cv::imshow("Result in RGB", h_result2);
  cv::imshow("Result in HSV ", h_result3);
  cv::imshow("Result in YCrCb ", h_result4);

  cv::waitKey();
  return 0;
}

imshow函数期望以 BGR 颜色格式传递彩色图像，因此使用imshow显示其他颜色格式的输出可能不会很吸引人。以下是在不同颜色格式下使用相同图像的前一个程序的输出：

图像阈值化

图像阈值化是一种非常简单的图像分割技术，用于根据某些强度值从灰度图像中提取重要区域。在这种技术中，如果像素值大于某个阈值值，则分配一个值，否则分配另一个值。

OpenCV 和 CUDA 中用于图像阈值化的函数是cv::cuda::threshold。此函数有许多参数。第一个参数是源图像，它应该是一个灰度图像。第二个参数是结果要存储的目标。第三个参数是阈值值，它用于分割像素值。第四个参数是maxVal常量，它表示如果像素值超过阈值值时赋予的值。OpenCV 提供不同类型的阈值化技术，由函数的最后一个参数决定。以下是一些阈值化类型：

cv:.THRESH_BINARY：如果像素的强度大于阈值，则将该像素强度设置为maxVal常量。否则将该像素强度设置为零。
cv::THRESH_BINARY_INV：如果像素的强度大于阈值，则将该像素强度设置为零。否则将该像素强度设置为maxVal常量。
cv::THRESH_TRUNC：这基本上是一个截断操作。如果像素的强度大于阈值，则将该像素强度设置为阈值。否则，保持强度值不变。
cv::THRESH_TOZERO：如果像素的强度大于阈值，则保持像素强度不变。否则将该像素强度设置为零。
cv::THRESH_TOZERO_INV：如果像素的强度大于阈值，则将该像素强度设置为零。否则保持像素强度不变。

实现所有这些阈值技术使用 OpenCV 和 CUDA 的程序如下：

#include <iostream>
#include "opencv2/opencv.hpp"

int main (int argc, char* argv[])
{
  cv::Mat h_img1 = cv::imread("images/cameraman.tif", 0);
  //Define device variables
  cv::cuda::GpuMat d_result1,d_result2,d_result3,d_result4,d_result5, d_img1;
  //Upload image on device
  d_img1.upload(h_img1);

  //Perform different thresholding techniques on device
  cv::cuda::threshold(d_img1, d_result1, 128.0, 255.0, cv::THRESH_BINARY);
  cv::cuda::threshold(d_img1, d_result2, 128.0, 255.0, cv::THRESH_BINARY_INV);
  cv::cuda::threshold(d_img1, d_result3, 128.0, 255.0, cv::THRESH_TRUNC);
  cv::cuda::threshold(d_img1, d_result4, 128.0, 255.0, cv::THRESH_TOZERO);
  cv::cuda::threshold(d_img1, d_result5, 128.0, 255.0, cv::THRESH_TOZERO_INV);

  cv::Mat h_result1,h_result2,h_result3,h_result4,h_result5;
  //Copy results back to host
  d_result1.download(h_result1);
  d_result2.download(h_result2);
  d_result3.download(h_result3);
  d_result4.download(h_result4);
  d_result5.download(h_result5);
  cv::imshow("Result Threshhold binary ", h_result1);
  cv::imshow("Result Threshhold binary inverse ", h_result2);
  cv::imshow("Result Threshhold truncated ", h_result3);
  cv::imshow("Result Threshhold truncated to zero ", h_result4);
  cv::imshow("Result Threshhold truncated to zero inverse ", h_result5);
  cv::waitKey();

  return 0;
}

在所有阈值技术的cv::cuda::threshold函数中，128 被用作像素强度的阈值，这是黑色（0）和白色（255）之间的中点。maxVal常量被设置为 255，当像素强度超过阈值时将用于更新像素强度。其他程序与之前看到的其他 OpenCV 程序类似。

程序的输出如下，显示了输入图像以及所有五种阈值技术的输出：

带有和没有 CUDA 支持的 OpenCV 应用程序的性能比较

可以通过处理单个图像所需的时间来衡量图像处理算法的性能。当算法在视频上工作时，性能是通过每秒帧数来衡量的，这表示它在一秒内可以处理的帧数。当算法每秒可以处理超过 30 帧时，它可以被认为是实时工作的。我们还可以衡量在 OpenCV 中实现的算法的性能，这将在本节中讨论。

如我们之前讨论的，当 OpenCV 与 CUDA 兼容性构建时，它可以显著提高算法的性能。OpenCV 在 CUDA 模块中的函数被优化以利用 GPU 并行处理能力。OpenCV 还提供了仅在 CPU 上运行的类似函数。在本节中，我们将比较上一节中构建的带有和不使用 GPU 的阈值操作的性能。我们将从处理一张图像所需的时间和每秒帧数来比较阈值操作的性能。以下是在 CPU 上实现阈值并测量性能的代码：

#include <iostream>
#include "opencv2/opencv.hpp"
using namespace cv;
using namespace std;

int main (int argc, char* argv[])
{
  cv::Mat src = cv::imread("images/cameraman.tif", 0);
  cv::Mat result_host1,result_host2,result_host3,result_host4,result_host5;

  //Get initial time in miliseconds
  int64 work_begin = getTickCount(); 
  cv::threshold(src, result_host1, 128.0, 255.0, cv::THRESH_BINARY);
  cv::threshold(src, result_host2, 128.0, 255.0,   cv::THRESH_BINARY_INV);
  cv::threshold(src, result_host3, 128.0, 255.0, cv::THRESH_TRUNC);
  cv::threshold(src, result_host4, 128.0, 255.0, cv::THRESH_TOZERO);
  cv::threshold(src, result_host5, 128.0, 255.0, cv::THRESH_TOZERO_INV);

  //Get time after work has finished     
  int64 delta = getTickCount() - work_begin;
  //Frequency of timer
  double freq = getTickFrequency();
  double work_fps = freq / delta;
  std::cout<<"Performance of Thresholding on CPU: " <<std::endl;
  std::cout <<"Time: " << (1/work_fps) <<std::endl;
  std::cout <<"FPS: " <<work_fps <<std::endl;
  return 0;
}

在前面的代码中，使用了来自 cv 命名空间的阈值函数，该函数仅使用 CPU 进行执行，而不是 cv::cuda 模块。算法的性能使用 gettickcount 和 gettickfrequency 函数进行测量。gettickcount 函数返回自系统启动后经过的毫秒数。我们测量了在图像操作代码执行前后的时间滴答。时间滴答之间的差异表示算法执行过程中处理图像所经过的滴答数。这个时间被测量在 delta 变量中。gettickfrequncy 函数返回计时器的频率。处理图像所需的总时间可以通过将时间滴答除以计时器频率来测量。这个时间的倒数表示 每秒帧数 （FPS）。这两个性能指标都打印在控制台上，用于 CPU 上的阈值应用。控制台输出如下：

如输出所示，CPU 处理一个图像需要 0.169766 秒，相当于 5.89046 FPS。现在我们将同样的算法实现到 GPU 上，并尝试测量代码的性能。根据之前的讨论，这应该会极大地提高算法的性能。GPU 实现的代码如下：

#include <iostream>
#include "opencv2/opencv.hpp"

int main (int argc, char* argv[])
{
  cv::Mat h_img1 = cv::imread("images/cameraman.tif", 0);
  cv::cuda::GpuMat d_result1,d_result2,d_result3,d_result4,d_result5, d_img1;
  //Measure initial time ticks
  int64 work_begin = getTickCount(); 
  d_img1.upload(h_img1);
  cv::cuda::threshold(d_img1, d_result1, 128.0, 255.0,   cv::THRESH_BINARY);
  cv::cuda::threshold(d_img1, d_result2, 128.0, 255.0,   cv::THRESH_BINARY_INV);
  cv::cuda::threshold(d_img1, d_result3, 128.0, 255.0, cv::THRESH_TRUNC);
  cv::cuda::threshold(d_img1, d_result4, 128.0, 255.0, cv::THRESH_TOZERO);
  cv::cuda::threshold(d_img1, d_result5, 128.0, 255.0, cv::THRESH_TOZERO_INV);

  cv::Mat h_result1,h_result2,h_result3,h_result4,h_result5;
  d_result1.download(h_result1);
  d_result2.download(h_result2);
  d_result3.download(h_result3);
  d_result4.download(h_result4);
  d_result5.download(h_result5);
  //Measure difference in time ticks
  int64 delta = getTickCount() - work_begin;
  double freq = getTickFrequency();
  //Measure frames per second
  double work_fps = freq / delta;
  std::cout <<"Performance of Thresholding on GPU: " <<std::endl;
  std::cout <<"Time: " << (1/work_fps) <<std::endl;
  std::cout <<"FPS: " <<work_fps <<std::endl;
  return 0;
}

在代码中，使用了来自 cv::cuda 模块的函数，该模块针对 GPU 并行处理能力进行了优化。图像被复制到设备内存，在 GPU 上进行操作，然后复制回主机。性能指标以类似的方式计算，并在控制台上打印。程序输出如下：

如所示，GPU 实现仅需要 0.55 毫秒来处理单个图像，相当于 1816 FPS。这比 CPU 实现有显著改进，尽管必须记住这是一个非常简单的应用，并不适合于 CPU 和 GPU 之间的性能比较。这个应用仅展示如何测量 OpenCV 中任何代码的性能。

通过运行 OpenCV 安装在 samples/gpu 目录中的示例代码，可以更真实地比较 CPU 和 GPU 的性能。其中一个代码，hog.cpp，从图像中计算 方向直方图 （HoG）特征，并使用 支持向量机 （SVM）进行分类。尽管算法的细节超出了本书的范围，但它给你一个关于使用 GPU 实现时性能改进的想法。在摄像头视频上的性能比较如下：

如所见，当我们仅使用 CPU 时，代码的性能大约为 13 FPS，如果我们使用 GPU，则性能提升至 24 FPS，这几乎是 CPU 性能的两倍。这会让你对使用 CUDA 与 OpenCV 结合的重要性有所了解。

总结来说，在本节中，我们比较了使用 CUDA（GPU）和不使用 CUDA（CPU）时 OpenCV 的性能。这再次强调了使用 CUDA 将极大地提高计算机视觉应用性能的观点。

摘要

在本章中，我们首先介绍了计算机视觉和图像处理。我们描述了 OpenCV 库，它专门为计算机视觉应用而设计，以及它与其他计算机视觉软件的不同之处。OpenCV 可以通过使用 CUDA 利用 GPU 的并行处理能力。我们查看在所有操作系统上安装具有 CUDA 的 OpenCV 的安装过程。我们描述了从磁盘读取图像、在屏幕上显示它以及将其保存回磁盘的过程。视频不过是图像的序列。我们学习了如何处理来自磁盘的视频以及从摄像头捕获的视频。我们开发了几个图像处理应用程序，对图像执行不同的操作，例如算术运算、逻辑运算、颜色空间转换和阈值处理。在最后一节中，我们比较了相同算法在 CPU 和 GPU 上的性能，包括处理图像所需的时间和 FPS。因此，在本章结束时，你对 OpenCV 与 CUDA 在计算机视觉应用中的有用性以及如何使用它编写简单代码有了了解。在下一章中，我们将在此基础上尝试开发更多有用的计算机视觉应用，例如使用 OpenCV 进行滤波、边缘检测和形态学操作。

问题

说明计算机视觉和图像处理这两个术语之间的区别
为什么 OpenCV 非常适合在嵌入式系统上部署计算机视觉应用
编写一个 OpenCV 命令，以红色初始化 1960 x 1960 彩色图像
编写一个程序，从网络摄像头捕获帧并将其保存到磁盘
OpenCV 用于读取和显示彩色图像的颜色格式是什么
编写一个程序，从网络摄像头捕获视频，将其转换为灰度并显示在屏幕上
编写一个程序来测量 GPU 上加法和减法操作的性能
编写一个程序进行图像的位与和或操作，并解释它如何用于遮罩

第六章：使用 OpenCV 和 CUDA 的基本计算机视觉操作

上一章描述了使用 OpenCV 和 CUDA 处理图像和视频的过程。我们查看了一些基本的图像和视频处理应用的代码，并比较了带有和没有 CUDA 加速的 OpenCV 代码的性能。在这一章中，我们将在此基础上构建知识，并尝试使用 OpenCV 和 CUDA 开发一些更多的计算机视觉和图像处理应用。这一章描述了在彩色和灰度图像中访问单个像素强度的方法。直方图是图像处理中的一个非常有用的概念。这一章描述了计算直方图的方法以及直方图均衡化如何提高图像的视觉效果。这一章还将描述如何使用 OpenCV 和 CUDA 执行不同的几何变换。图像滤波是一个非常重要的概念，它在图像预处理和特征提取中非常有用。这一章将详细描述这一点。本章的最后部分将描述不同的形态学操作，如腐蚀、膨胀、开运算和闭运算。

本章将涵盖以下主题：

在 OpenCV 中访问单个像素强度
直方图计算和直方图均衡化
图像变换
图像的滤波操作
图像的形态学操作

技术要求

本章需要具备图像处理和计算机视觉的基本理解。它需要熟悉基本的 C 或 C++ 编程语言、CUDA 以及前几章中解释的所有示例代码。本章中使用的所有代码都可以从以下 GitHub 链接下载：github.com/PacktPublishing/Hands-On-GPU-Accelerated-Computer-Vision-with-OpenCV-and-CUDA。代码可以在任何操作系统上执行，尽管它只在 Ubuntu 16.04 上进行了测试。

查看以下视频以查看代码的实际效果：

bit.ly/2xERUDL

访问图像的单个像素强度

当我们处理图像时，有时需要访问特定位置的像素强度值。当我们想要改变一组像素的亮度或对比度，或者想要执行其他像素级操作时，这非常有用。对于一个 8 位灰度图像，该点的强度值将在 0 到 255 的范围内，而对于彩色图像，将会有三个不同的强度值，分别对应蓝色、绿色和红色通道，所有这些通道的值都在 0 到 255 之间。

OpenCV 提供了一个cv::Mat::at<>方法来访问任何通道图像在特定位置的强度值。它需要一个参数，即要访问强度值的位置。该点通过Point类传递，行和列值作为参数。对于灰度图像，该方法将返回一个标量对象，而对于彩色图像，它将返回一个包含三个强度的向量。访问灰度图像以及彩色图像在特定位置的像素强度的代码如下：

#include <iostream>
#include "opencv2/opencv.hpp"
int main ()
{
  //Gray Scale Image
  cv::Mat h_img1 = cv::imread("images/cameraman.tif",0);
  cv::Scalar intensity = h_img1.at<uchar>(cv::Point(100, 50));
  std::cout<<"Pixel Intensity of gray scale Image at (100,50) is:"  <<intensity.val[0]<<std::endl;
  //Color Image
  cv::Mat h_img2 = cv::imread("images/autumn.tif",1);
  cv::Vec3b intensity1 = h_img1.at<cv::Vec3b>(cv::Point(100, 50));
  std::cout<<"Pixel Intensity of color Image at (100,50) is:"<<intensity1<<std::endl;
  return 0;
}

灰度图像首先被读取，然后在这个图像对象上调用at方法。强度值在(100,50)点被测量，这表示第 100 行和第 50 列的像素。它返回一个标量，存储在强度变量中。该值被打印在控制台上。对于彩色图像，遵循相同的程序，但返回值将是一个包含三个强度的向量，存储在Vec3b对象中。强度值被打印在控制台上。上述程序的输出如下：

如所示，灰度图像在(100,50)处的像素强度为9，而对于彩色图像则是[175,179,177]，这表示蓝色强度为175，绿色强度为179，红色强度为177。同样的方法用于修改特定位置的像素强度。假设你想将(100,50)位置的像素强度改为128，则可以编写如下代码：

h_img1.at<uchar>(100, 50) = 128;

总结来说，在本节中我们看到了一种访问和改变特定位置强度值的方法。在下一节中，我们将看到在 OpenCV 中计算直方图的方法。

OpenCV 中的直方图计算和平滑

直方图是图像的一个重要属性，因为它提供了该图像外观的全局描述。可以从直方图中获得大量信息。它代表了图像中灰度级别的相对出现频率。它基本上是在 X 轴上灰度级别和 Y 轴上每个灰度级别的像素数的图表。如果直方图集中在左侧，则图像会非常暗；如果集中在右侧，则图像会非常亮。为了获得良好的图像视觉质量，它应该均匀分布。

以下图像展示了暗、亮和正常图像的直方图：

OpenCV 提供了一个函数来计算图像的直方图。该函数的语法如下：

void cv::cuda::calcHist ( InputArray src, OutputArray hist)

函数需要两个数组作为参数。第一个数组是需要计算直方图的输入图像。第二个参数是输出数组，其中将存储直方图。输出可以绘制成直方图，如图像前面的截图所示。如前所述，平坦的直方图可以提高图像的视觉质量。OpenCV 和 CUDA 提供了一个函数来平坦直方图，这在下一节中将有描述。

直方图均衡化

完美的图像在其所有灰度级别中像素数量相等。因此，直方图应该具有大的动态范围和整个范围内的像素数量相等。这可以通过一种称为直方图均衡化的技术来实现。它是任何计算机视觉应用中非常重要的预处理步骤。在本节中，我们将看到如何使用 OpenCV 和 CUDA 对灰度图像和彩色图像进行直方图均衡化。

灰度图像

灰度图像通常是 8 位单通道图像，具有 256 个不同的灰度级别。如果直方图分布不均匀，图像可能太暗或太亮，此时应进行直方图均衡化以改善图像的视觉质量。以下代码描述了在灰度图像上进行直方图均衡化的过程：

#include <iostream>
#include "opencv2/opencv.hpp"
int main ()
{
  cv::Mat h_img1 = cv::imread("images/cameraman.tif",0);
  cv::cuda::GpuMat d_img1,d_result1;
  d_img1.upload(h_img1);
  cv::cuda::equalizeHist(d_img1, d_result1);
  cv::Mat h_result1;
  d_result1.download(h_result1);
  cv::imshow("Original Image ", h_img1);
  cv::imshow("Histogram Equalized Image", h_result1);
  cv::waitKey();
  return 0;
}

读取的图像被上传到设备内存以进行直方图均衡化。这是一个计算密集型的步骤，因此 CUDA 加速将有助于提高程序的性能。OpenCV 提供了equalizeHist函数用于直方图均衡化。它需要两个参数。第一个参数是源图像，第二个参数是目标图像。目标图像被下载回主机并在控制台上显示。直方图均衡化后的输出如下：

如所见，直方图均衡化后的图像在视觉质量上优于原始图像。接下来将描述对彩色图像进行相同操作的过程。

彩色图像

直方图均衡化也可以应用于彩色图像。它必须在单独的通道上执行。因此，彩色图像必须分成三个通道。每个通道的直方图独立均衡，然后合并通道以重建图像。以下是对彩色图像进行直方图均衡化的代码：

#include <iostream>
#include "opencv2/opencv.hpp"
int main ()
{
  cv::Mat h_img1 = cv::imread("images/autumn.tif");
  cv::Mat h_img2,h_result1;
  cvtColor(h_img1, h_img2, cv::COLOR_BGR2HSV);
  //Split the image into 3 channels; H, S and V channels respectively and store it in a std::vector
  std::vector< cv::Mat > vec_channels;
  cv::split(h_img2, vec_channels); 
  //Equalize the histogram of only the V channel 
  cv::equalizeHist(vec_channels[2], vec_channels[2]);
  //Merge 3 channels in the vector to form the color image in HSV color space.
  cv::merge(vec_channels, h_img2); 
  //Convert the histogram equalized image from HSV to BGR color space again
  cv::cvtColor(h_img2,h_result1, cv::COLOR_HSV2BGR);
  cv::imshow("Original Image ", h_img1);
  cv::imshow("Histogram Equalized Image", h_result1);
  cv::waitKey();
  return 0;
}

在 BGR 颜色空间中，直方图通常不进行均衡化；使用 HSV 和 YCrCb 颜色空间进行均衡化。因此，在代码中，将 BGR 颜色空间转换为 HSV 颜色空间。然后，使用split函数将其拆分为三个独立的通道。现在，色调和饱和度通道包含颜色信息，因此没有必要均衡这些通道。直方图均衡化仅在值通道上执行。使用merge函数将三个通道合并回重建彩色图像。使用imshow将 HSV 彩色图像转换回 BGR 颜色空间进行显示。程序的输出如下：

总结来说，直方图均衡化可以提升图像的视觉效果，因此它是任何计算机视觉应用中非常重要的预处理步骤。下一节将描述图像的几何变换。

图像的几何变换

有时，在更大的计算机视觉应用中，需要缩放图像、平移图像和旋转图像。本节将解释这类几何变换。

图像缩放

在某些计算机视觉应用中，图像需要具有特定的尺寸。因此，需要将任意大小的图像转换为特定尺寸。OpenCV 提供了一个用于缩放图像的函数。图像缩放的代码如下：

#include <iostream>
#include "opencv2/opencv.hpp"
int main ()
{
  cv::Mat h_img1 = cv::imread("images/cameraman.tif",0);
  cv::cuda::GpuMat d_img1,d_result1,d_result2;
  d_img1.upload(h_img1);
  int width= d_img1.cols;
  int height = d_img1.size().height;
  cv::cuda::resize(d_img1,d_result1,cv::Size(200, 200),   cv::INTER_CUBIC);
  cv::cuda::resize(d_img1,d_result2,cv::Size(0.5*width, 0.5*height),   cv::INTER_LINEAR); 
  cv::Mat h_result1,h_result2;
  d_result1.download(h_result1);
  d_result2.download(h_result2);
  cv::imshow("Original Image ", h_img1);
  cv::imshow("Resized Image", h_result1);
  cv::imshow("Resized Image 2", h_result2);
  cv::waitKey();
  return 0;
}

可以使用两个不同的函数获取图像的高度和宽度，如代码所示。Mat对象的rows和cols属性分别描述图像的height和width。Mat对象还有一个size()方法，它具有height和width属性，用于查找图像的大小。图像以两种方式缩放。在第一种方式中，图像被缩放为特定的(200,200)大小，在第二种方式中，它被缩放为其原始尺寸的一半。OpenCV 提供了resize函数来执行此操作。它有四个参数。

前两个参数分别是源图像和目标图像。第三个参数是目标图像的大小。它使用Size对象定义。当图像缩放时，必须从源图像对目标图像上的像素值进行插值。有各种插值方法可供使用，例如双线性插值、双三次插值和面积插值。这些插值方法作为resize函数的第四个参数提供。它可以设置为cv::INTER_LINEAR (双线性)、cv::INTER_CUBIC (双三次)或cv::INTER_AREA (面积)。图像缩放的代码输出如下：

图像平移和旋转

图像平移和旋转是某些计算机视觉应用中需要的重要几何变换。OpenCV 提供了一个简单的 API 来在图像上执行这些变换。执行平移和旋转的代码如下：

#include <iostream>
#include "opencv2/opencv.hpp"

int main ()
{
  cv::Mat h_img1 = cv::imread("images/cameraman.tif",0);
  cv::cuda::GpuMat d_img1,d_result1,d_result2;
  d_img1.upload(h_img1);
  int cols= d_img1.cols;
  int rows = d_img1.size().height;
  //Translation
  cv::Mat trans_mat = (cv::Mat_<double>(2,3) << 1, 0, 70, 0, 1, 50);
  cv::cuda::warpAffine(d_img1,d_result1,trans_mat,d_img1.size());
  //Rotation
  cv::Point2f pt(d_img1.cols/2., d_img1.rows/2.); 
  cv::Mat rot_mat = cv::getRotationMatrix2D(pt, 45, 1.0);
  cv::cuda::warpAffine(d_img1, d_result2, rot_mat, cv::Size(d_img1.cols, d_img1.rows));
  cv::Mat h_result1,h_result2;
  d_result1.download(h_result1);
  d_result2.download(h_result2);
  cv::imshow("Original Image ", h_img1);
  cv::imshow("Translated Image", h_result1);
  cv::imshow("Rotated Image", h_result2);
  cv::waitKey();
  return 0;
}

需要创建一个平移矩阵，该矩阵指定了图像在水平和垂直方向上的平移。它是一个 2 x 3 的矩阵，如下所示：

tx 和 ty 是沿 x 和 y 方向的平移偏移。在代码中，这个矩阵使用 Mat 对象创建，其中 X-方向的偏移量为 70，Y-方向的偏移量为 50。这个矩阵作为参数传递给 warpAffine 函数以实现图像平移。warpAffine 函数的其他参数分别是源图像、目标图像和输出图像的大小。

应该创建一个旋转矩阵，用于在特定点以特定角度旋转图像。OpenCV 提供了 cv::getRotationMatrix2D 函数来构建这个旋转矩阵。它需要三个参数。第一个参数是旋转点；在这种情况下使用图像的中心。第二个参数是旋转角度，指定为 45 度。最后一个参数是缩放比例，指定为 1。构建的旋转矩阵再次作为参数传递给 warpAffine 函数以实现图像旋转。

图像平移和图像旋转代码的输出如下：

总结来说，本节描述了使用 OpenCV 和 CUDA 实现的各种几何变换，如图像缩放、图像平移和图像旋转。

图像上的滤波操作

到目前为止描述的方法都是针对单个像素强度进行的，被称为点处理方法。有时查看像素的邻域而不是仅查看单个像素强度是有帮助的。这些被称为邻域处理技术。邻域可以是 3 x 3、5 x 5、7 x 7 等等，并且以特定像素为中心。图像滤波是邻域处理中的一个重要技术。

过滤是信号处理中的一个重要概念，其中我们拒绝一定频率范围的信号，并允许一定频率范围的信号通过。图像中是如何测量频率的？如果一个区域的灰度值变化缓慢，那么它是一个低频区域。如果灰度值变化剧烈，那么它是一个高频区域。通常，图像的背景被认为是低频区域，而边缘是高频区域。卷积是邻域处理和图像滤波中的一个非常重要的数学概念。它将在下一节中解释。

图像上的卷积操作

卷积的基本思想源于生物学中类似的概念，称为感受野，其中对图像中的某些部分敏感，而对其他部分不敏感。它可以用以下方式数学表示：

g(x,y)=f(x,y)h(x,y)= ∑∑f(n,m)h(x-n,y-m)*

简化形式下，此方程是滤波器h与图像的子图像f（以(x,y)点为中心）的点积。此乘积的答案是图像中的(x,y)点g。为了说明卷积操作在图像上的工作原理，以下图显示了将 3 x 3 滤波器应用于 6 x 6 大小图像的示例：

通过将最左侧的红色窗口与滤波器进行点积，以找到目标图像中的一个点。点积的答案将是 2 ((11 + 11 + 10 + 15 + 11 +17 +10 +11 + 11)/9)*。将此窗口向右移动 1 个像素后，重复相同的操作，答案将是 3。这将对图像中的所有窗口重复进行，以构建目标图像。通过改变 3 x 3 滤波器矩阵的值，可以构建不同的低通和高通滤波器。这将在下一两节中解释。

图像上的低通滤波

低通滤波器从图像中移除高频内容。通常，噪声被认为是高频内容，因此低通滤波器从图像中移除噪声。有许多类型的噪声，如高斯噪声、均匀噪声、指数噪声和盐和胡椒噪声，这些都会影响图像。低通滤波器用于消除这类噪声。有许多类型的低通滤波器可用：

平均或箱式滤波器
高斯滤波器
中值滤波器

本节解释了这些滤波器及其使用 OpenCV 的实现。

平均滤波器

平均滤波器，正如其名所示，对邻域像素执行平均操作。如果图像中存在高斯噪声，则可以使用低通平均滤波器来去除噪声。由于平均操作，它还会模糊图像的边缘。邻域可以是 3 x 3、5 x 5、7 x 7 等。滤波器窗口的尺寸越大，图像的模糊程度就越高。3 x 3 和 5 x 5 平均掩模如下：

OpenCV 提供了一个简单的接口，可以在图像上应用许多类型的滤波器。以下代码演示了如何使用不同掩模应用平均滤波器：

#include <iostream>
#include "opencv2/opencv.hpp"
int main ()
{
  cv::Mat h_img1 = cv::imread("images/cameraman.tif",0);
  cv::cuda::GpuMat d_img1,d_result3x3,d_result5x5,d_result7x7;
  d_img1.upload(h_img1);
  cv::Ptr<cv::cuda::Filter> filter3x3,filter5x5,filter7x7;
  filter3x3 = cv::cuda::createBoxFilter(CV_8UC1,CV_8UC1,cv::Size(3,3));
  filter3x3->apply(d_img1, d_result3x3);
  filter5x5 = cv::cuda::createBoxFilter(CV_8UC1,CV_8UC1,cv::Size(5,5));
  filter5x5->apply(d_img1, d_result5x5);
  filter7x7 = cv::cuda::createBoxFilter(CV_8UC1,CV_8UC1,cv::Size(7,7));
  filter7x7->apply(d_img1, d_result7x7);

  cv::Mat h_result3x3,h_result5x5,h_result7x7;
  d_result3x3.download(h_result3x3);
  d_result5x5.download(h_result5x5);
  d_result7x7.download(h_result7x7);
  cv::imshow("Original Image ", h_img1);
  cv::imshow("Blurred with kernel size 3x3", h_result3x3);
  cv::imshow("Blurred with kernel size 5x5", h_result5x5);
  cv::imshow("Blurred with kernel size 7x7", h_result7x7);
  cv::waitKey();
  return 0;
}

cv::Ptr是一个用于智能指针的模板类，用于存储cv::cuda::Filter类型的过滤器。然后，使用createBoxFilter函数创建不同窗口大小的平均滤波器。它需要三个必选参数和三个可选参数。第一个和第二个参数是源图像和目标图像的数据类型。它们被假定为CV_8UC1，表示 8 位无符号灰度图像。第三个参数定义了滤波器窗口的大小。它可以是一个 3 x 3、5 x 5、7 x 7 等。第四个参数是锚点，其默认值为(-1,-1)，表示锚点位于核的中心点。最后两个可选参数与像素插值方法和边界值相关，这里省略。

创建的过滤器指针有一个应用方法，该方法用于将创建的过滤器应用于任何图像。它有三个参数。第一个参数是源图像，第二个参数是目标图像，第三个可选参数是 CUDA 流，它用于多任务处理，如本书前面所述。在代码中，对图像应用了不同大小的三个平均滤波器。结果如下：

从输出中可以看出，随着滤波器大小的增加，用于平均的像素更多，这会在图像上引入更多的模糊。尽管大滤波器可以消除更多的噪声。

高斯滤波器

高斯滤波器使用具有高斯分布的掩码来过滤图像，而不是简单的平均掩码。此滤波器还引入了图像上的平滑模糊，并且广泛用于从图像中消除噪声。一个 5 x 5 的高斯滤波器，其标准差约为 1，如下所示：

OpenCV 提供了一个实现高斯滤波器的函数。其代码如下：

#include <iostream>
#include "opencv2/opencv.hpp"

int main ()
{
  cv::Mat h_img1 = cv::imread("images/cameraman.tif",0);
  cv::cuda::GpuMat d_img1,d_result3x3,d_result5x5,d_result7x7;
  d_img1.upload(h_img1);
  cv::Ptr<cv::cuda::Filter> filter3x3,filter5x5,filter7x7;
  filter3x3 = cv::cuda::createGaussianFilter(CV_8UC1,CV_8UC1,cv::Size(3,3),1);
  filter3x3->apply(d_img1, d_result3x3);
  filter5x5 = cv::cuda::createGaussianFilter(CV_8UC1,CV_8UC1,cv::Size(5,5),1);
  filter5x5->apply(d_img1, d_result5x5);
  filter7x7 = cv::cuda::createGaussianFilter(CV_8UC1,CV_8UC1,cv::Size(7,7),1);
  filter7x7->apply(d_img1, d_result7x7);

  cv::Mat h_result3x3,h_result5x5,h_result7x7;
  d_result3x3.download(h_result3x3);
  d_result5x5.download(h_result5x5);
  d_result7x7.download(h_result7x7);
  cv::imshow("Original Image ", h_img1);
  cv::imshow("Blurred with kernel size 3x3", h_result3x3);
  cv::imshow("Blurred with kernel size 5x5", h_result5x5);
  cv::imshow("Blurred with kernel size 7x7", h_result7x7);
  cv::waitKey();
  return 0;
}

createGaussianFilter函数用于创建高斯滤波器的掩码。源图像和目标图像的数据类型、滤波器大小以及水平方向的标准差作为参数提供给函数。我们还可以提供一个垂直方向的标准差作为参数；如果没有提供，则其默认值等于水平方向的标准差。使用apply方法将不同大小的创建的高斯掩码应用于图像。程序输出如下：

同样，随着高斯滤波器大小的增加，图像中引入了更多的模糊。高斯滤波器用于消除噪声并在图像上引入平滑的模糊。

中值滤波

当图像受到椒盐噪声的影响时，它不会被平均或高斯滤波器消除。它需要一个非线性滤波器。在邻域中进行中值运算而不是平均可以帮助消除椒盐噪声。在这个滤波器中，邻域中 9 个像素值的中值放置在中心像素上。它将消除由椒盐噪声引入的极端高或低值。尽管 OpenCV 和 CUDA 提供了一个中值滤波函数，但它的速度比 OpenCV 中的常规函数慢，因此使用以下代码实现中值滤波：

#include <iostream>
#include "opencv2/opencv.hpp"

int main ()
{
  cv::Mat h_img1 = cv::imread("images/saltpepper.png",0);
  cv::Mat h_result;
  cv::medianBlur(h_img1,h_result,3);
  cv::imshow("Original Image ", h_img1);
  cv::imshow("Median Blur Result", h_result);
  cv::waitKey();
  return 0;
}

OpenCV 中的medianBlur函数用于实现中值滤波。它需要三个参数。第一个参数是源图像，第二个参数是目标图像，第三个参数是中值操作的窗口大小。中值滤波的输出如下：

源图像受到椒盐噪声的影响，如图中的截图所示。这种噪声通过 3x3 大小的中值滤波器完全消除，而没有引入极端的模糊。因此，中值滤波是图像应用受到椒盐噪声影响时的一个非常重要的预处理步骤。

总结一下，我们看到了三种类型的低通滤波器，它们在各种计算机视觉应用中得到了广泛使用。平均滤波器和高斯滤波器用于消除高斯噪声，但它们也会模糊图像的边缘。中值滤波器用于去除椒盐噪声。

图像的高通滤波

高通滤波器从图像中移除低频分量并增强高频分量。因此，当高通滤波器应用于图像时，它会移除背景，因为它是低频区域，并增强边缘，这些是高频分量。因此，高通滤波器也可以称为边缘检测器。滤波器的系数将改变，否则它与上一节中看到的滤波器相似。有许多高通滤波器可用，如下所示：

Sobel 滤波器
Scharr 滤波器
拉普拉斯滤波器

在本节中，我们将分别看到它们中的每一个。

Sobel 滤波器

Sobel 算子或 Sobel 滤波器是一种广泛用于边缘检测应用的图像处理和计算机视觉算法。它是一个 3 x 3 的滤波器，用于近似图像强度函数的梯度。它提供了一个单独的滤波器来计算水平和垂直方向的梯度。该滤波器以与本章前面所述类似的方式与图像卷积。水平和垂直的 3 x 3 Sobel 滤波器如下：

实现此 Sobel 滤波器的代码如下：

#include <iostream>
#include "opencv2/opencv.hpp"

int main ()
{
  cv::Mat h_img1 = cv::imread("images/blobs.png",0);
  cv::cuda::GpuMat d_img1,d_resultx,d_resulty,d_resultxy;
  d_img1.upload(h_img1);
  cv::Ptr<cv::cuda::Filter> filterx,filtery,filterxy;
  filterx = cv::cuda::createSobelFilter(CV_8UC1,CV_8UC1,1,0);
  filterx->apply(d_img1, d_resultx);
  filtery = cv::cuda::createSobelFilter(CV_8UC1,CV_8UC1,0,1);
  filtery->apply(d_img1, d_resulty);
  cv::cuda::add(d_resultx,d_resulty,d_resultxy); 
  cv::Mat h_resultx,h_resulty,h_resultxy;
  d_resultx.download(h_resultx);
  d_resulty.download(h_resulty);
  d_resultxy.download(h_resultxy);
  cv::imshow("Original Image ", h_img1);
  cv::imshow("Sobel-x derivative", h_resultx);
  cv::imshow("Sobel-y derivative", h_resulty);
  cv::imshow("Sobel-xy derivative", h_resultxy);
  cv::waitKey();
  return 0;
}

OpenCV 提供了createSobelFilter函数来实现 Sobel 滤波器。它需要许多参数。前两个参数是源图像和目标图像的数据类型。第三个和第四个参数分别是x和y导数的阶数。对于计算x导数或垂直边缘，提供 1 和 0，而对于计算y导数或水平边缘，提供 0 和 1。第五个参数表示核的大小，是可选的。默认值是 3。也可以提供导数的比例。

要同时看到水平和垂直边缘，需要将x导数和y导数的结果相加。结果如下：

Sobel 算子提供了非常不精确的导数近似，但仍然在计算机视觉应用中的边缘检测方面非常有用。它没有旋转对称性；为了克服这一点，使用了 Scharr 算子。

Scharr 滤波器

由于 Sobel 没有提供旋转对称性，因此使用不同的滤波器掩码来克服这一点，如下所示：

从掩码中可以看出，Scharr 算子更重视中间行或中间列以找到边缘。实现 Scharr 滤波器的程序如下：

#include <iostream>
#include "opencv2/opencv.hpp"
int main ()
{
  cv::Mat h_img1 = cv::imread("images/blobs.png",0);
  cv::cuda::GpuMat d_img1,d_resultx,d_resulty,d_resultxy;
  d_img1.upload(h_img1);
  cv::Ptr<cv::cuda::Filter> filterx,filtery;
  filterx = cv::cuda::createScharrFilter(CV_8UC1,CV_8UC1,1,0);
  filterx->apply(d_img1, d_resultx);
  filtery = cv::cuda::createScharrFilter(CV_8UC1,CV_8UC1,0,1);
  filtery->apply(d_img1, d_resulty);
  cv::cuda::add(d_resultx,d_resulty,d_resultxy); 
  cv::Mat h_resultx,h_resulty,h_resultxy;
  d_resultx.download(h_resultx);
  d_resulty.download(h_resulty);
  d_resultxy.download(h_resultxy);
  cv::imshow("Original Image ", h_img1);
  cv::imshow("Scharr-x derivative", h_resultx);
  cv::imshow("Scharr-y derivative", h_resulty);
  cv::imshow("Scharr-xy derivative", h_resultxy);
  cv::waitKey();
    return 0;
}

OpenCV 提供了createScharrFilter函数来实现 Scharr 滤波器。它需要许多参数。前两个参数是源图像和目标图像的数据类型。第三个和第四个参数分别是x和y导数的阶数。对于计算x导数或垂直边缘，提供 1 和 0，而对于计算y导数或水平边缘，提供 0 和 1。第五个参数，表示核的大小，是可选的。默认值是 3。

要同时看到水平和垂直边缘，需要将x导数和y导数的结果相加。结果如下：

Laplacian 滤波器

Laplacian 滤波器也是一种导数算子，用于在图像中找到边缘。不同之处在于 Sobel 和 Scharr 是一阶导数算子，而 Laplacian 是二阶导数算子。它也同时找到水平和垂直方向的边缘，这与 Sobel 和 Scharr 算子不同。Laplacian 滤波器计算二阶导数，因此对图像中的噪声非常敏感，在应用 Laplacian 滤波器之前最好对图像进行模糊处理以去除噪声。实现 Laplacian 滤波器的代码如下：

#include <iostream>
#include "opencv2/opencv.hpp"

int main ()
{
  cv::Mat h_img1 = cv::imread("images/blobs.png",0);
  cv::cuda::GpuMat d_img1,d_result1,d_result3;
  d_img1.upload(h_img1);
  cv::Ptr<cv::cuda::Filter> filter1,filter3;
  filter1 = cv::cuda::createLaplacianFilter(CV_8UC1,CV_8UC1,1);
  filter1->apply(d_img1, d_result1);
  filter3 = cv::cuda::createLaplacianFilter(CV_8UC1,CV_8UC1,3);
  filter3->apply(d_img1, d_result3);
  cv::Mat h_result1,h_result3;
  d_result1.download(h_result1);
  d_result3.download(h_result3);
  cv::imshow("Original Image ", h_img1);
  cv::imshow("Laplacian filter 1", h_result1);
  cv::imshow("Laplacian filter 3", h_result3);
  cv::waitKey();
  return 0;
}

使用createLaplacianFilter函数在图像上应用了两个核大小为 1 和 3 的 Laplacian 滤波器。除了核的大小外，该函数还需要作为参数提供源图像和目标图像的数据类型。创建的 Laplacian 滤波器通过apply方法应用于图像。Laplacian 滤波器的输出如下：

总结来说，在本节中，我们描述了不同的高通滤波器，如 Sobel、Scharr 和 Laplacian 滤波器。Sobel 和 Scharr 是一阶导数算子，用于计算边缘，它们对噪声的敏感性较低。Laplacian 是一个二阶导数算子，用于计算边缘，它对噪声非常敏感。

图像的形态学操作

图像形态学处理图像的区域和形状。它用于提取对表示形状和区域有用的图像成分。图像形态学将图像视为集合的总和，这与之前看到的其他图像处理操作不同。图像与一个小模板相互作用，该模板称为结构元素，它定义了图像形态学中的感兴趣区域或邻域。本节中解释了可以在图像上执行的各种形态学操作，将逐一进行说明。

腐蚀：腐蚀将中心像素设置为邻域内所有像素的最小值。邻域由结构元素定义，它是一个由 1s 和 0s 组成的矩阵。腐蚀用于扩大物体中的孔洞，缩小边界，消除岛屿，并去除可能存在于图像边界上的狭窄半岛。
膨胀：膨胀将中心像素设置为邻域内所有像素的最大值。膨胀增加了白色块的大小，减少了黑色区域的大小。它用于填充物体中的孔洞并扩展物体的边界。
开运算：图像开运算基本上是腐蚀和膨胀的组合。图像开运算定义为腐蚀后跟膨胀。这两个操作都使用相同的结构元素执行。它用于平滑图像的轮廓，破坏狭窄的桥梁并隔离相互接触的物体。它在分析发动机油中的磨损颗粒、回收纸中的墨水颗粒等方面得到应用。
闭运算：图像闭运算定义为膨胀后跟腐蚀。这两个操作都使用相同的结构元素执行。它用于融合狭窄的裂缝并消除小孔。

形态学算子可以通过将它们应用于仅包含黑白两色的二值图像来轻松理解。OpenCV 和 CUDA 提供了一个简单的 API 来在图像上应用形态学变换。相应的代码如下：

#include <iostream>
#include "opencv2/opencv.hpp"
int main ()
{
  cv::Mat h_img1 = cv::imread("images/blobs.png",0);
  cv::cuda::GpuMat d_img1,d_resulte,d_resultd,d_resulto, d_resultc;
  cv::Mat element = cv::getStructuringElement(cv::MORPH_RECT,cv::Size(5,5)); 
  d_img1.upload(h_img1);
  cv::Ptr<cv::cuda::Filter> filtere,filterd,filtero,filterc;
  filtere = cv::cuda::createMorphologyFilter(cv::MORPH_ERODE,CV_8UC1,element);
  filtere->apply(d_img1, d_resulte);
  filterd = cv::cuda::createMorphologyFilter(cv::MORPH_DILATE,CV_8UC1,element);
  filterd->apply(d_img1, d_resultd);
  filtero = cv::cuda::createMorphologyFilter(cv::MORPH_OPEN,CV_8UC1,element);
  filtero->apply(d_img1, d_resulto);
  filterc = cv::cuda::createMorphologyFilter(cv::MORPH_CLOSE,CV_8UC1,element);
  filterc->apply(d_img1, d_resultc);

  cv::Mat h_resulte,h_resultd,h_resulto,h_resultc;
  d_resulte.download(h_resulte);
  d_resultd.download(h_resultd);
  d_resulto.download(h_resulto);
  d_resultc.download(h_resultc);
  cv::imshow("Original Image ", h_img1);
  cv::imshow("Erosion", h_resulte);
  cv::imshow("Dilation", h_resultd);
  cv::imshow("Opening", h_resulto);
  cv::imshow("closing", h_resultc);
  cv::waitKey();
  return 0;
}

首先需要创建一个定义形态学操作邻域的结构元素。这可以通过使用 OpenCV 中的getStructuringElement函数来完成。需要将结构元素的形状和尺寸作为参数传递给此函数。在代码中，定义了一个 5 x 5 大小的矩形结构元素。

形态学操作的过滤器是通过使用createMorphologyFilter函数创建的。它需要三个必填参数。第一个参数定义要执行的操作。cv::MORPH_ERODE用于腐蚀，cv::MORPH_DILATE用于膨胀，cv::MORPH_OPEN用于开运算，cv::MORPH_CLOSE用于闭运算。第二个参数是图像的数据类型，第三个参数是之前创建的结构元素。使用apply方法将这些过滤器应用于图像。

图像上形态学操作的输出如下：

从输出结果可以看出，腐蚀操作会减小物体的边界，而膨胀操作则会使其变厚。我们将白色部分视为物体，黑色部分则是背景。开运算可以平滑图像的轮廓。闭运算可以消除图像中的小孔。如果将结构元素的尺寸从 5 x 5 增加到 7 x 7，那么在腐蚀操作中边界腐蚀会更加明显，而在膨胀操作中边界会变得更厚。在 5 x 5 腐蚀图像中可见的左侧的小圆圈，在用 7 x 7 尺寸腐蚀时会被移除。

使用 7 x 7 结构元素的形态学操作的输出如下：

总结来说，形态学操作对于找出定义图像形状和区域的组件非常重要。它可以用来填充图像中的孔洞并平滑图像的轮廓。

摘要

本章介绍了在图像中特定位置访问像素强度的方法。当我们对图像进行逐点操作时，这非常有用。直方图是描述图像的一个非常重要的全局特征。本章介绍了计算直方图和直方图均衡化的方法，这可以提高图像的视觉效果。详细解释了各种几何变换，如图像缩放、旋转和平移。图像滤波是一种有用的邻域处理技术，用于消除噪声和提取图像的边缘特征，并进行了详细描述。低通滤波器用于去除噪声，但它也会模糊图像的边缘。高通滤波器去除背景，这是一个低频区域，同时增强边缘，这些是高频区域。本章的最后部分介绍了不同的形态学操作，如腐蚀、膨胀、开运算和闭运算，这些可以用来描述图像的形状并填充图像中的空洞。在下一章中，我们将使用这些概念，结合 OpenCV 和 CUDA 构建一些有用的计算机视觉应用。

问题

编写一个 OpenCV 函数，在控制台打印任何彩色图像在位置（200,200）的像素强度。
编写一个 OpenCV 函数，将图像调整到（300,200）像素大小。使用双线性插值方法。
编写一个 OpenCV 函数，通过 2 倍上采样图像。使用面积插值方法。
判断对错：随着平均滤波器大小的增加，模糊程度降低。
判断对错：中值滤波器可以去除高斯噪声。
可以采取哪些步骤来降低拉普拉斯算子的噪声敏感性？
编写一个 OpenCV 函数来实现顶帽和黑帽形态学操作。

第七章：使用 OpenCV 和 CUDA 进行目标检测和跟踪

上一章描述了使用 OpenCV 和 CUDA 的基本计算机视觉操作。在这一章中，我们将看到如何使用这些基本操作以及 OpenCV 和 CUDA 来开发复杂的计算机视觉应用。我们将使用目标检测和跟踪的例子来展示这个概念。目标检测和跟踪是计算机视觉中一个非常活跃的研究领域。它涉及在图像中识别物体的位置并在一系列帧中跟踪它。基于颜色、形状和图像的其他显著特征，已经提出了许多用于此任务的方法。在这一章中，这些算法使用 OpenCV 和 CUDA 实现。我们首先解释基于颜色的物体检测，然后描述检测特定形状物体的方法。所有物体都有显著的特性，可以用来检测和跟踪物体。本章描述了不同特征检测算法的实现以及如何使用它们来检测物体。本章的最后部分将演示使用背景减除技术，该技术将前景与背景分离以进行目标检测和跟踪。

本章将涵盖以下主题：

目标检测和跟踪简介
基于颜色的目标检测和跟踪
基于形状的目标检测和跟踪
基于特征的物体检测
使用 Haar 级联的目标检测
背景减除方法

技术要求

本章需要具备良好的图像处理和计算机视觉理解。它还需要一些关于用于目标检测和跟踪的算法的基本知识。需要熟悉基本的 C 或 C++编程语言、CUDA 以及前几章中解释的所有代码。本章中使用的所有代码都可以从以下 GitHub 链接下载：github.com/PacktPublishing/Hands-On-GPU-Accelerated-Computer-Vision-with-OpenCV-and-CUDA。代码可以在任何操作系统上执行，尽管它只在 Ubuntu 16.04 上进行了测试。

查看以下视频以查看代码的实际应用：

bit.ly/2PSRqkU

目标检测和跟踪简介

目标检测和跟踪是计算机视觉领域的一个活跃的研究课题，它通过一系列帧努力检测、识别和跟踪物体。已经发现，视频序列中的目标检测和跟踪是一个具有挑战性的任务，并且是一个非常耗时的过程。目标检测是构建更大计算机视觉系统的第一步。可以从检测到的物体中推导出大量信息，如下所示：

检测到的目标可以被分类到特定的类别
可以在图像序列中进行跟踪
可以从检测到的对象中获取更多关于场景或其他对象推断的信息

目标跟踪被定义为在视频的每一帧中检测对象，并建立从一帧到另一帧检测到的对象的对应关系。

目标检测与跟踪的应用

目标检测与跟踪可用于开发视频监控系统以跟踪可疑活动、事件和人员。它可以用于开发智能交通系统以跟踪车辆和检测交通违规行为。在自动驾驶汽车中，目标检测对于提供周围环境信息和规划导航至关重要。它对于自动驾驶员辅助系统中的行人检测或车辆检测也非常有用。它可用于医疗领域的应用，如乳腺癌检测或脑瘤检测等。它可用于面部和手势识别。它在工业装配和生产线的质量控制中具有广泛的应用。对于搜索引擎中的图像检索和照片管理也非常重要。

目标检测的挑战

目标检测是一个具有挑战性的任务，因为现实生活中的图像会受到噪声、光照变化、动态背景、阴影效果、相机抖动和运动模糊的影响。当要检测的对象旋转、缩放或被遮挡时，目标检测变得困难。许多应用需要检测多个对象类别。如果检测的类别数量很大，那么处理速度就成为一个重要问题，同时系统在处理这些类别时，如何不损失准确性也是一个关键问题。

有许多算法可以克服这些挑战中的一些。这些算法在本章中进行了讨论。本章没有详细描述这些算法，但更侧重于如何使用 CUDA 和 OpenCV 来实现它们。

基于颜色的目标检测与跟踪

一个对象有许多全局特征，如颜色和形状，这些特征描述了对象的整体。这些特征可以用于在一系列帧中检测对象并跟踪它。在本节中，我们将使用颜色作为特征来检测具有特定颜色的对象。当要检测的对象是特定颜色且与背景颜色不同时，这种方法很有用。如果对象和背景颜色相同，则这种检测方法将失败。在本节中，我们将尝试使用 OpenCV 和 CUDA 从网络摄像头流中检测任何蓝色对象。

蓝色目标检测与跟踪

第一个问题可能是应该使用哪种颜色空间来分割蓝色。红绿蓝（RGB）颜色空间没有将颜色信息与强度信息分开。将颜色信息与强度信息分开的颜色空间，如色调饱和度值（HSV）和YCrCb（其中 Y′是亮度分量，CB 和 CR 是蓝差和红差色度分量），对于这类任务非常理想。每种颜色在色调通道中都有一个特定的范围，可以用来检测该颜色。以下是从开始摄像头、捕获帧到上传设备内存以进行 GPU 操作的样板代码：

#include <iostream>
#include "opencv2/opencv.hpp"

using namespace cv;
using namespace std;

int main()
{
  VideoCapture cap(0); //capture the video from web cam
  // if webcam is not available then exit the program
  if ( !cap.isOpened() ) 
  {
    cout << "Cannot open the web cam" << endl;
    return -1;
  }
  while (true)
  {
    Mat frame;
    // read a new frame from webcam
    bool flag = cap.read(frame); 
    if (!flag) 
    {
      cout << "Cannot read a frame from webcam" << endl;
      break;
    }

    cuda::GpuMat d_frame, d_frame_hsv,d_intermediate,d_result;
    cuda::GpuMat d_frame_shsv[3];
    cuda::GpuMat d_thresc[3];
    Mat h_result;
    d_frame.upload(frame);

    d_result.download(h_result);
    imshow("Thresholded Image", h_result); 
    imshow("Original", frame); 

    if (waitKey(1) == 'q') 
    {
      break; 
    }
  }
  return 0;
}
}

要检测蓝色，我们需要在 HSV 颜色空间中找到蓝色颜色的范围。如果范围准确，则检测将准确。蓝色颜色在三个通道（色调、饱和度和值）中的范围如下：

lower_range = [110,50,50]
upper_range = [130,255,255]

这个范围将被用来在特定通道中对图像进行阈值处理，以创建蓝色颜色的掩码。如果这个掩码再次与原始帧进行 AND 操作，那么结果图像中就只会剩下蓝色对象。以下是这个操作的代码：

//Transform image to HSV
cuda::cvtColor(d_frame, d_frame_hsv, COLOR_BGR2HSV);

//Split HSV 3 channels
cuda::split(d_frame_hsv, d_frame_shsv);

//Threshold HSV channels for blue color according to range
cuda::threshold(d_frame_shsv[0], d_thresc[0], 110, 130, THRESH_BINARY);
cuda::threshold(d_frame_shsv[1], d_thresc[1], 50, 255, THRESH_BINARY);
cuda::threshold(d_frame_shsv[2], d_thresc[2], 50, 255, THRESH_BINARY);

//Bitwise AND the channels
cv::cuda::bitwise_and(d_thresc[0], d_thresc[1],d_intermediate);
cv::cuda::bitwise_and(d_intermediate, d_thresc[2], d_result);

摄像头捕获的帧被转换为 HSV 颜色空间。蓝色在这三个通道中具有不同的范围，因此每个通道都需要单独进行阈值处理。使用split方法将通道分割，并使用threshold函数进行阈值处理。每个通道的最小和最大范围用作下限和上限阈值。在此范围内的通道值将被转换为白色，其他则转换为黑色。这三个阈值通道通过逻辑 AND 操作得到一个用于蓝色颜色的最终掩码。这个掩码可以用来从视频中检测和跟踪具有蓝色颜色的对象。

两个帧的输出，一个没有蓝色对象，另一个有蓝色对象，如下所示：

从结果可以看出，当帧中不包含任何蓝色对象时，掩码几乎为黑色；而在下面的帧中，当蓝色对象进入画面时，该部分变为白色。这种方法仅在背景不包含对象颜色时才会有效。

基于形状的对象检测与跟踪

对象的形状也可以作为全局特征来检测具有独特形状的对象。这个形状可以是直线、多边形、圆形或任何其他不规则形状。对象边界、边缘和轮廓可以用来检测具有特定形状的对象。在本节中，我们将使用 Canny 边缘检测算法和 Hough 变换来检测两个规则形状，即直线和圆形。

Canny 边缘检测

在上一章中，我们看到了各种高通滤波器，这些滤波器可以用作边缘检测器。在本节中，我们使用 OpenCV 和 CUDA 实现了结合高斯滤波、梯度查找、非极大值抑制和阈值滞后的 Canny 边缘检测算法。正如上一章所解释的，高通滤波器对噪声非常敏感。在 Canny 边缘检测中，在检测边缘之前先进行高斯平滑，这使得它对噪声的敏感性降低。它还在检测边缘后有一个非极大值抑制阶段，以从结果中去除不必要的边缘。

Canny 边缘检测是一个计算密集型任务，难以用于实时应用。算法的 CUDA 版本可以用来加速它。实现 Canny 边缘检测算法的代码描述如下：

#include <cmath>
#include <iostream>
#include "opencv2/opencv.hpp"

using namespace std;
using namespace cv;
using namespace cv::cuda;

int main()
{
  Mat h_image = imread("images/drawing.JPG",0);
  if (h_image.empty())
  {
    cout << "can not open image"<< endl;
    return -1;
  }
  GpuMat d_edge,d_image;
  Mat h_edge;
  d_image.upload(h_image);
  cv::Ptr<cv::cuda::CannyEdgeDetector> Canny_edge = cv::cuda::createCannyEdgeDetector(2.0, 100.0, 3, false);
  Canny_edge->detect(d_image, d_edge);
  d_edge.download(h_edge);
  imshow("source", h_image);
  imshow("detected edges", h_edge);
  waitKey(0);

  return 0;
}

OpenCV 和 CUDA 提供了createCannyEdgeDetector类用于 Canny 边缘检测。创建此类的对象时，可以传递许多参数。前两个参数是阈值滞后的低阈值和高阈值。如果某一点的强度梯度大于最大阈值，则将其分类为边缘点。如果梯度小于低阈值，则该点不是边缘点。如果梯度在阈值之间，则根据连通性决定该点是否为边缘点。第三个参数是边缘检测器的孔径大小。最后一个参数是布尔参数，表示是否使用L2_norm或L1_norm进行梯度幅度计算。L2_norm计算成本较高，但更准确。真值表示使用L2_norm。代码的输出如下所示：

您可以调整下限和上限阈值，以更准确地检测给定图像的边缘。边缘检测是许多计算机视觉应用的重要预处理步骤，Canny 边缘检测被广泛用于此目的。

使用霍夫变换进行直线检测

在许多计算机视觉应用中，如车道检测，检测直线非常重要。它还可以用来检测其他规则形状的一部分直线。霍夫变换是计算机视觉中用于检测直线的流行特征提取技术。我们不会详细介绍霍夫变换如何检测直线，但我们将看到它如何在 OpenCV 和 CUDA 中实现。实现霍夫变换进行直线检测的代码如下：

#include <cmath>
#include <iostream>
#include "opencv2/opencv.hpp"

using namespace std;
using namespace cv;
using namespace cv::cuda;

int main()
{
  Mat h_image = imread("images/drawing.JPG",0);
  if (h_image.empty())
  {
    cout << "can not open image"<< endl;
    return -1;
  }

  Mat h_edge;
  cv::Canny(h_image, h_edge, 100, 200, 3);

  Mat h_imagec;
  cv::cvtColor(h_edge, h_imagec, COLOR_GRAY2BGR);
  Mat h_imageg = h_imagec.clone();
  GpuMat d_edge, d_lines;
  d_edge.upload(h_edge);
  {
    const int64 start = getTickCount();
    Ptr<cuda::HoughSegmentDetector> hough = cuda::createHoughSegmentDetector(1.0f, (float) (CV_PI / 180.0f), 50, 5);
    hough->detect(d_edge, d_lines);

    const double time_elapsed = (getTickCount() - start) / getTickFrequency();
    cout << "GPU Time : " << time_elapsed * 1000 << " ms" << endl;
    cout << "GPU FPS : " << (1/time_elapsed) << endl;
  }
  vector<Vec4i> lines_g;
  if (!d_lines.empty())
  {
    lines_g.resize(d_lines.cols);
    Mat h_lines(1, d_lines.cols, CV_32SC4, &lines_g[0]);
    d_lines.download(h_lines);
  }
  for (size_t i = 0; i < lines_g.size(); ++i)
  {
    Vec4i line_point = lines_g[i];
    line(h_imageg, Point(line_point[0], line_point[1]), Point(line_point[2], line_point[3]), Scalar(0, 0, 255), 2, LINE_AA);
  }

  imshow("source", h_image);
  imshow("detected lines [GPU]", h_imageg);
  waitKey(0);
  return 0;
}

OpenCV 提供了createHoughSegmentDetector类用于实现霍夫变换。它需要一个图像的边缘图作为输入。因此，使用 Canny 边缘检测器从图像中检测边缘。Canny 边缘检测器的输出上传到设备内存以进行 GPU 计算。正如上一节所讨论的，边缘也可以在 GPU 上计算。

创建了createHoughSegmentDetector对象。它需要许多参数。第一个参数表示在霍夫变换中使用的参数r的分辨率，通常取 1 像素。第二个参数是参数 theta 的弧度分辨率，通常取 1 弧度或π/180。第三个参数是需要形成线条的最小点数，通常取 50 像素。最后一个参数是考虑为同一直线的两个点之间的最大间隔，通常取 5 像素。

创建的对象的检测方法用于检测直线。它需要两个参数。第一个参数是要检测边缘的图像，第二个参数是存储检测到的线条点的数组。该数组包含检测到的线条的起始和结束(x,y)点。使用 OpenCV 的线条函数通过for循环迭代该数组，在图像上绘制单个线条。最终图像使用imshow函数显示。

霍夫变换是一个数学密集型步骤。为了展示 CUDA 的优势，我们将实现 CPU 和 CUDA 的相同算法，并比较它们的性能。以下是对 CPU 霍夫变换的代码：

Mat h_imagec; 
vector<Vec4i> h_lines;
{
  const int64 start = getTickCount();
  HoughLinesP(h_edge, h_lines, 1, CV_PI / 180, 50, 60, 5);
  const double time_elapsed = (getTickCount() - start) / getTickFrequency();
  cout << "CPU Time : " << time_elapsed * 1000 << " ms" << endl;
  cout << "CPU FPS : " << (1/time_elapsed) << endl;
}

for (size_t i = 0; i < h_lines.size(); ++i)
{
  Vec4i line_point = h_lines[i];
  line(h_imagec, Point(line_point[0], line_point[1]), Point(line_point[2], line_point[3]), Scalar(0, 0, 255), 2, LINE_AA);
}
imshow("detected lines [CPU]", h_imagec);

使用HoughLinesP函数在 CPU 上通过概率霍夫变换检测线条。前两个参数是源图像和存储输出线条点的数组。第三个和第四个参数是r和 theta 的分辨率。第五个参数是表示线条的最小交点数的阈值。第六个参数表示形成线条所需的最小点数。最后一个参数表示考虑在相同线条上的点的最大间隔。

函数返回的数组使用for循环迭代，以在原始图像上显示检测到的线条。GPU 和 CPU 函数的输出如下：

以下截图显示了 GPU 和 CPU 代码在霍夫变换性能上的比较：

在 CPU 上处理单个图像大约需要 4 毫秒，在 GPU 上需要 1.5 毫秒，这在 CPU 上相当于 248 FPS，在 GPU 上相当于 632 FPS，几乎是 GPU 上 2.5 倍的提升。

圆检测

霍夫变换也可以用于圆检测。它可以用于许多应用，如球检测和跟踪以及硬币检测等，在这些应用中，对象是圆形的。OpenCV 和 CUDA 提供了一个类来实现这一点。以下是用霍夫变换进行硬币检测的代码：

#include "opencv2/opencv.hpp"
#include <iostream>

using namespace cv;
using namespace std;

int main(int argc, char** argv)
{
  Mat h_image = imread("images/eight.tif", IMREAD_COLOR);
  Mat h_gray;
  cvtColor(h_image, h_gray, COLOR_BGR2GRAY);
  cuda::GpuMat d_gray,d_result;
  std::vector<cv::Vec3f> d_Circles;
cv::Ptr<cv::cuda::HoughCirclesDetector> detector = cv::cuda::createHoughCirclesDetector(1, 100, 122, 50, 1, max(h_image.size().width, h_image.size().height));
  d_gray.upload(h_gray);
  detector->detect(d_gray, d_result);
  d_Circles.resize(d_result.size().width);
  if (!d_Circles.empty())
    d_result.row(0).download(cv::Mat(d_Circles).reshape(3, 1));

  cout<<"No of circles: " <<d_Circles.size() <<endl;
  for( size_t i = 0; i < d_Circles.size(); i++ )
  {
    Vec3i cir = d_Circles[i];
    circle( h_image, Point(cir[0], cir[1]), cir[2], Scalar(255,0,0), 2, LINE_AA);
  }
  imshow("detected circles", h_image);
  waitKey(0);

  return 0;
}

有一个createHoughCirclesDetector类用于检测圆形物体。创建了该类的对象。在创建该类的对象时可以提供许多参数。第一个参数是dp，表示累加器分辨率与图像分辨率的倒数，通常取为 1。第二个参数是检测到的圆心之间的最小距离。第三个参数是 Canny 阈值，第四个参数是累加器阈值。第五和第六个参数是要检测的圆的最小和最大半径。

圆心之间的最小距离取为100像素。您可以尝试调整这个值。如果这个值减小，那么原始图像上会错误地检测到许多圆，而如果这个值增加，那么一些真正的圆可能会被错过。最后两个参数，即最小和最大半径，如果不知道确切的尺寸，可以取为0。在前面的代码中，它被设置为1和图像的最大尺寸，以检测图像中的所有圆。程序输出如下：

Hough 变换对高斯噪声和椒盐噪声非常敏感。因此，在应用 Hough 变换之前，有时最好先使用高斯和中值滤波器对图像进行预处理。这将给出更准确的结果。

总结来说，我们使用了 Hough 线变换和圆变换来检测具有规则形状的物体。轮廓和凸性也可以用于形状检测。这些功能在 OpenCV 中可用，但 CUDA 实现中不可用。您将不得不开发这些函数的自己的版本。

关键点检测器和描述符

到目前为止，我们使用了全局特征，如颜色和形状来检测物体。这些特征易于计算，快速，并且需要的内存量小，但它们只能在已有关于物体的某些信息的情况下使用。如果没有这些信息，则使用局部特征，这些特征需要更多的计算和内存，但它们更准确。在本节中，解释了寻找局部特征的多种算法。它们也被称为关键点检测器。关键点是表征图像的点，可以用来精确地定义一个物体。

来自加速区域测试（FAST）特征检测器的特征

FAST 算法用于从图像中检测角点作为关键点。它通过对每个像素应用段测试来检测角点。它考虑像素周围的 16 像素圆。如果在半径为 16 的圆中有连续的n个点，其强度大于Ip+t或小于Ip-t，则该像素被认为是角点。Ip是像素p的强度，t是选定的阈值。

有时，不是检查半径内的所有点，而是检查几个选定的点以确定强度值来决定角落点。这加速了 FAST 算法的性能。FAST 提供了可以作为关键点利用的角落点来检测对象。它是旋转不变的，因为即使对象旋转，对象的角落也会保持不变。FAST 不是尺度不变的，因为尺寸的增加可能会导致强度值的平滑过渡，而不是在角落处的尖锐过渡。

OpenCV 和 CUDA 提供了一个高效实现 FAST 算法的方法。以下是用 FAST 算法检测关键点的程序：

#include <iostream>
#include "opencv2/opencv.hpp"

using namespace cv;
using namespace std;

int main()
{
  Mat h_image = imread( "images/drawing.JPG", 0 );

  //Detect the key-points using FAST Detector
  cv::Ptr<cv::cuda::FastFeatureDetector> detector = cv::cuda::FastFeatureDetector::create(100,true,2);
  std::vector<cv::key point> key-points;
  cv::cuda::GpuMat d_image;
  d_image.upload(h_image);
  detector->detect(d_image, key-points);
  cv::drawkey-points(h_image,key-points,h_image);
  //Show detected key-points
  imshow("Final Result", h_image );
  waitKey(0);
  return 0;
}

OpenCV 和 CUDA 提供了一个FastFeatureDetector类来实现 FAST 算法。这个类的对象是通过类的 create 方法创建的。它需要三个参数。第一个参数是要用于 FAST 算法的强度阈值。第二个参数指定是否使用非最大抑制。它是一个布尔值，可以指定为true或false。第三个参数表示用于计算邻域的 FAST 方法。有三个方法，cv2.FAST_FEATURE_DETECTOR_TYPE_5_8、cv2.FAST_FEATURE_DETECTOR_TYPE_7_12和cv2.FAST_FEATURE_DETECTOR_TYPE_9_16，可以作为标志0、1或2指定。

创建的对象的 detect 方法用于检测关键点。它需要一个输入图像和一个存储关键点的向量作为参数。可以使用drawkey-points函数在原始图像上绘制计算出的关键点。它需要源图像、关键点的向量和目标图像作为参数。

可以改变强度阈值以检测不同数量的关键点。如果阈值低，则更多的关键点会通过分割测试并被归类为关键点。随着阈值的增加，检测到的关键点数量将逐渐减少。同样，如果非最大抑制是假的，则在单个角落点可能会检测到多个关键点。以下是代码的输出：

从输出结果可以看出，随着阈值从 10 增加到 50 和 100，关键点的数量减少。这些关键点可以用于检测查询图像中的对象。

定向 FAST 和旋转 BRIEF (ORB) 特征检测

ORB 是一个非常高效的特徵检测和描述算法。它是特征检测的 FAST 算法和特征描述的二进制鲁棒独立基本特征（BRIEF）算法的组合。它为广泛用于对象检测的 SURF 和 SIFT 算法提供了一个高效的替代方案。由于它们是专利的，使用它们需要付费。ORB 在无需付费的情况下匹配 SIFT 和 SURF 的性能。

OpenCV 和 CUDA 提供了一个易于实现的 ORB 算法的 API。实现 ORB 算法的代码如下：

#include <iostream>
#include "opencv2/opencv.hpp"

using namespace cv;
using namespace std;

int main()
{
  Mat h_image = imread( "images/drawing.JPG", 0 );
  cv::Ptr<cv::cuda::ORB> detector = cv::cuda::ORB::create();
  std::vector<cv::key point> key-points;
  cv::cuda::GpuMat d_image;
  d_image.upload(h_image);
  detector->detect(d_image, key-points);
  cv::drawkey-points(h_image,key-points,h_image);
  imshow("Final Result", h_image );
  waitKey(0);
  return 0;
}

ORB类的对象是通过create方法创建的。此方法的所有参数都是可选的，因此我们使用了它的默认值。创建的对象的detect方法用于从图像中检测关键点。它需要一个输入图像和关键点向量的向量，输出将存储在这些参数中。使用drawkey-points函数在图像上绘制检测到的关键点。前述代码的输出如下：

ORB类还提供了一个方法来计算所有关键点的描述符。这些描述符可以准确地描述对象，并可用于从图像中检测对象。这些描述符也可以用于对对象进行分类。

加速鲁棒特征检测和匹配

SURF 通过基于简单二维盒滤波器的计算来近似高斯拉普拉斯。使用积分图像可以轻松计算与盒滤波器的卷积，这提高了算法的性能。SURF 依赖于 Hessian 矩阵的行列式来处理尺度和位置。Hessian 矩阵的近似行列式可以表示为：

其中，w是滤波器响应的相对权重，用于平衡行列式的表达式。Dx、Dy是拉普拉斯算子在X和Y方向的结果。

SURF 在水平和垂直方向上使用小波响应，使用积分图像方法进行方向分配。还应用了适当的 Gaussian 权重。通过计算角度为 60 度的滑动方向窗口内所有响应的总和来估计主导方向。

对于特征描述，SURF 在水平和垂直方向上使用 Haar 小波响应。这是对图像中的所有子区域进行计算的结果，从而得到一个具有 64 个维度的 SURF 特征描述符。维度越低，计算和匹配的速度越快。为了提高精度，SURF 特征描述符还有一个扩展的 128 维版本。SURF 是旋转不变和尺度不变的。

与使用 128 维特征向量的 SIFT 相比，SURF 具有更高的处理速度，因为它使用 64 方向的特性向量。SURF 擅长处理模糊和旋转的图像，但不擅长处理视点和光照变化。

OpenCV 和 CUDA 提供了一个 API 来计算 SURF 关键点和描述符。我们还将看到如何使用这些 API 在查询图像中检测对象。SURF 特征检测和匹配的代码如下：

#include <stdio.h>
#include <iostream>
#include "opencv2/opencv.hpp"
#include "opencv2/features2d.hpp"
#include "opencv2/xfeatures2d.hpp"
#include "opencv2/xfeatures2d/nonfree.hpp"
#include "opencv2/xfeatures2d/cuda.hpp"

using namespace cv;
using namespace cv::xfeatures2d;
using namespace std;

int main( int argc, char** argv )
{
  Mat h_object_image = imread( "images/object1.jpg", 0 ); 
  Mat h_scene_image = imread( "images/scene1.jpg", 0 );
  cuda::GpuMat d_object_image;
  cuda::GpuMat d_scene_image;
  cuda::GpuMat d_key-points_scene, d_key-points_object; 
  vector< key point > h_key-points_scene, h_key-points_object;
  cuda::GpuMat d_descriptors_scene, d_descriptors_object;
  d_object_image.upload(h_object_image);
  d_scene_image.upload(h_scene_image);
  cuda::SURF_CUDA surf(150);
  surf( d_object_image, cuda::GpuMat(), d_key-points_object, d_descriptors_object );
surf( d_scene_image, cuda::GpuMat(), d_key-points_scene, d_descriptors_scene );

Ptr< cuda::DescriptorMatcher > matcher = cuda::DescriptorMatcher::createBFMatcher();
vector< vector< DMatch> > d_matches;
matcher->knnMatch(d_descriptors_object, d_descriptors_scene, d_matches, 3);
surf.downloadkey-points(d_key-points_scene, h_key-points_scene);
surf.downloadkey-points(d_key-points_object, h_key-points_object);
std::vector< DMatch > good_matches;
for (int k = 0; k < std::min(h_key-points_object.size()-1, d_matches.size()); k++)
{
  if ( (d_matches[k][0].distance < 0.75*(d_matches[k][1].distance)) &&
      ((int)d_matches[k].size() <= 2 && (int)d_matches[k].size()>0) )
  {
    good_matches.push_back(d_matches[k][0]);
  }
}
std::cout << "size:" <<good_matches.size();
Mat h_image_result;
drawMatches( h_object_image, h_key-points_object, h_scene_image, h_key-points_scene,
      good_matches, h_image_result, Scalar::all(-1), Scalar::all(-1),
      vector<char>(), DrawMatchesFlags::DEFAULT );
imshow("Good Matches & Object detection", h_image_result);
waitKey(0);
return 0;
}

从磁盘读取两张图像。第一张图像包含要检测的对象。第二张图像是要搜索对象的查询图像。我们将从这两张图像中计算 SURF 特征，然后匹配这些特征以从查询图像中检测对象。

OpenCV 提供了用于计算 SURF 特征的 SURF_CUDA 类。该类的对象被创建。它需要一个 Hessian 阈值作为参数。这里取值为 150。这个阈值决定了 Hessian 行列式计算输出的点必须有多大，才能被认为是关键点。更大的阈值值将导致更少但更显著的兴趣点，而较小的值将导致更多但不太显著的点。可以根据应用来选择。

这个 surf 对象用于从对象和查询图像中计算关键点和描述符。图像、图像的数据类型、存储关键点的向量以及描述符作为参数传递。为了匹配查询图像中的对象，需要匹配两张图像中的描述符。OpenCV 提供了不同的匹配算法来实现这个目的，如 Brute-Force 匹配器和 快速近似最近邻库（FLANN）匹配器。

程序中使用的是 Brute-Force 匹配器；这是一个简单的方法。它使用某种距离计算方法，将对象中与查询图像中所有其他特征匹配的特征描述符取出来。它返回最佳匹配关键点，或使用最近邻算法通过 matcher 类的 knnMatch 方法返回最佳 k 个匹配。knnMatch 方法需要两组描述符以及最近邻的数量。在代码中取值为 3。

从 knnMatch 方法返回的匹配点中提取出良好的匹配关键点。这些良好的匹配是通过使用原始论文中描述的比率测试方法找到的。这些良好的匹配用于从场景中检测对象。

使用 drawMatches 函数在两张图像的匹配良好点之间画线。它需要许多参数。第一个参数是源图像，第二个参数是源图像的关键点，第三个参数是第二张图像，第四个参数是第二张图像的关键点，第五个参数是输出图像。第六个参数是线条和关键点的颜色。这里取值为 Scalar::all(-1)，表示将随机选择颜色。第七个参数是关键点的颜色，这些关键点没有匹配。它也取值为 Scalar::all(-1)，表示将随机选择颜色。最后两个参数指定了绘制匹配的掩码和标志设置。使用空掩码，以便绘制所有匹配。

这些匹配可以用来在检测到的物体周围绘制边界框，这将定位场景中的物体。绘制边界框的代码如下：

std::vector<Point2f> object;
std::vector<Point2f> scene;
for (int i = 0; i < good_matches.size(); i++) {
  object.push_back(h_key-points_object[good_matches[i].queryIdx].pt);
  scene.push_back(h_key-points_scene[good_matches[i].trainIdx].pt);
}
Mat Homo = findHomography(object, scene, RANSAC);
std::vector<Point2f> corners(4);
std::vector<Point2f> scene_corners(4);
corners[0] = Point(0, 0);
corners[1] = Point(h_object_image.cols, 0);
corners[2] = Point(h_object_image.cols, h_object_image.rows);
corners[3] = Point(0, h_object_image.rows);
perspectiveTransform(corners, scene_corners, Homo);
line(h_image_result, scene_corners[0] + Point2f(h_object_image.cols, 0),scene_corners[1] + Point2f(h_object_image.cols, 0), Scalar(255, 0, 0), 4);
line(h_image_result, scene_corners[1] + Point2f(h_object_image.cols, 0),scene_corners[2] + Point2f(h_object_image.cols, 0),Scalar(255, 0, 0), 4);
line(h_image_result, scene_corners[2] + Point2f(h_object_image.cols, 0),scene_corners[3] + Point2f(h_object_image.cols, 0),Scalar(255, 0, 0), 4);
line(h_image_result, scene_corners[3] + Point2f(h_object_image.cols, 0),scene_corners[0] + Point2f(h_object_image.cols, 0),Scalar(255, 0, 0), 4);

OpenCV 提供了findHomography函数，用于根据良好的匹配搜索场景中物体的位置、方向和比例。前两个参数是从物体和场景图像中提取的良好匹配的关键点。随机样本一致性（RANSAC）方法作为参数之一传递，用于找到最佳平移矩阵。

找到这个平移矩阵后，使用perspectiveTransform函数来找到物体。它需要一个四角点和平移矩阵作为参数。这些变换点用于在检测到的物体周围绘制边界框。用于查找特征和匹配物体的 SURF 程序的输出如下：

该图包含物体图像、查询图像和检测图像。从前面的图像可以看出，SURF 可以准确地确定物体的位置，即使物体被旋转。尽管有时它可能会检测到错误特征。可以通过改变 Hessian 阈值和测试比率来找到最佳匹配。

因此，总结一下，在本节中，我们看到了 FAST、ORB 和 SURF 关键点检测算法。我们还看到了如何使用这些点通过使用 SURF 特征作为示例来匹配和定位图像中的物体。您也可以尝试使用 FAST 和 ORB 特征来完成相同的工作。在下一节中，我们将详细讨论用于从图像中检测面部和眼睛的 Haar 级联。

使用 Haar 级联进行物体检测

Haar 级联使用矩形特征来检测物体。它使用不同大小的矩形来计算不同的线和边缘特征。矩形包含一些黑白区域，如图所示，它们在图像中的不同位置居中：

Haar-like 特征选择算法背后的思想是计算矩形内部白色像素总和与黑色像素总和之间的差异。

这种方法的主要优势是使用积分图像快速进行求和计算。这使得 Haar 级联非常适合实时物体检测。它处理图像所需的时间比之前描述的 SURF 等算法少。由于计算量较小且内存占用较少，该算法也可以在嵌入式系统（如 Raspberry Pi）上实现。它被称为 Haar-like，因为它基于与 Haar 小波相同的原理。Haar 级联在人体检测中广泛使用，包括面部和眼部检测等部分。它还可以用于表情分析。Haar 级联可用于检测车辆等物体。

在本节中，描述了使用 Haar 级联从图像和摄像头中检测人脸和眼睛的方法。Haar 级联是一种机器学习算法，需要对其进行训练以执行特定任务。对于特定应用从头开始训练 Haar 级联是困难的，因此 OpenCV 提供了一些训练好的 XML 文件，可用于检测对象。这些 XML 文件位于 OpenCV 或 CUDA 安装的opencv\data\haarcascades_cuda文件夹中。

使用 Haar 级联进行人脸检测

在本节中，我们将使用 Haar 级联从图像和实时摄像头中检测人脸。使用 Haar 级联从图像中检测人脸的代码如下：

#include "opencv2/objdetect/objdetect.hpp"
#include "opencv2/highgui/highgui.hpp"
#include "opencv2/imgproc/imgproc.hpp"
#include "opencv2/cudaobjdetect.hpp" 
#include <iostream>
#include <stdio.h>

using namespace std;
using namespace cv;

int main( )
{
  Mat h_image;
  h_image = imread("images/lena_color_512.tif", 0); 
  Ptr<cuda::CascadeClassifier> cascade = cuda::CascadeClassifier::create("haarcascade_frontalface_alt2.xml");
  cuda::GpuMat d_image;
  cuda::GpuMat d_buf;
  d_image.upload(h_image);
  cascade->detectMultiScale(d_image, d_buf);
  std::vector<Rect> detections;
  cascade->convert(d_buf, detections);
  if (detections.empty())
    std::cout << "No detection." << std::endl;
  cvtColor(h_image,h_image,COLOR_GRAY2BGR);
  for(int i = 0; i < detections.size(); ++i)
  {
    rectangle(h_image, detections[i], Scalar(0,255,255), 5);
  }
  imshow("Result image", h_image);
  waitKey(0); 
  return 0;
}

OpenCV 和 CUDA 提供了CascadeClassifier类，可用于实现 Haar 级联。使用 create 方法创建该类的对象。它需要加载训练好的 XML 文件的文件名。创建的对象有detectMultiScale方法，可以从图像中检测到多个尺度的对象。它需要一个图像文件和一个Gpumat数组作为参数来存储输出结果。使用CascadeClassifier对象的 convert 方法将此gpumat向量转换为标准矩形向量。此转换向量包含绘制检测到的对象矩形坐标。

detectMultiScale函数有许多参数可以在调用函数之前修改。这些包括用于指定每次图像缩放时图像大小将减少多少的scaleFactor，以及指定每个矩形应保留的最小邻居数minNeighbors；minSize指定最小对象大小，maxSize指定最大对象大小。所有这些参数都有默认值，所以在正常情况下通常不需要修改。如果我们想更改它们，那么在调用detectMultiscale函数之前可以使用以下代码：

cascade->setMinNeighbors(0);
cascade->setScaleFactor(1.01);

第一个函数将设置最小邻居数为0，第二个函数将在每次缩放后通过一个因子1.01减小图像大小。缩放因子对于检测不同大小物体的检测非常重要。如果它很大，则算法完成所需时间会更短，但可能有些面部无法检测到。如果它很小，则算法完成所需时间会更长，并且会更准确。前述代码的输出如下：

从视频中

Haar 级联的相同概念可以用于从视频中检测人脸。检测人脸的代码包含在while循环中，以便在视频的每一帧中检测到人脸。从摄像头进行人脸检测的代码如下：

#include <iostream>
#include <opencv2/opencv.hpp>
using namespace cv;
using namespace std;

int main()
{
  VideoCapture cap(0);
  if (!cap.isOpened()) {
    cerr << "Can not open video source";
    return -1;
  }
  std::vector<cv::Rect> h_found;
  cv::Ptr<cv::cuda::CascadeClassifier> cascade = cv::cuda::CascadeClassifier::create("haarcascade_frontalface_alt2.xml");
  cv::cuda::GpuMat d_frame, d_gray, d_found;
  while(1)
  {
    Mat frame;
    if ( !cap.read(frame) ) {
      cerr << "Can not read frame from webcam";
      return -1;
    }
    d_frame.upload(frame);
    cv::cuda::cvtColor(d_frame, d_gray, cv::COLOR_BGR2GRAY);

    cascade->detectMultiScale(d_gray, d_found);
    cascade->convert(d_found, h_found);

    for(int i = 0; i < h_found.size(); ++i)
    {
      rectangle(frame, h_found[i], Scalar(0,255,255), 5);
    }

    imshow("Result", frame);
    if (waitKey(1) == 'q') {
      break;
    }
  }

  return 0;
}

初始化摄像头并逐个捕获摄像头帧。此帧上传到设备内存中进行 GPU 处理。通过使用类的create方法创建CascadeClassifier类的对象。创建对象时提供面部检测的 XML 文件作为参数。在while循环内部，对每个帧应用detectMultiscale方法，以便在每帧中检测不同大小的面部。使用convert方法将检测到的位置转换为矩形向量。然后使用for循环迭代此向量，以便使用rectangle函数在所有检测到的面部上绘制边界框。程序输出如下：

使用 Haar 级联进行眼睛检测

本节将描述在检测人类眼睛中使用 Haar 级联的方法。用于眼睛检测的已训练 Haar 级联的 XML 文件位于 OpenCV 安装目录中。此文件用于检测眼睛。其代码如下：

#include <iostream>
#include <stdio.h>
 #include <opencv2/opencv.hpp>

using namespace std;
using namespace cv;

int main( )
{
  Mat h_image;
  h_image = imread("images/lena_color_512.tif", 0); 
  Ptr<cuda::CascadeClassifier> cascade = cuda::CascadeClassifier::create("haarcascade_eye.xml");
  cuda::GpuMat d_image;
  cuda::GpuMat d_buf;
  d_image.upload(h_image);
  cascade->setScaleFactor(1.02);
  cascade->detectMultiScale(d_image, d_buf);
  std::vector<Rect> detections;
  cascade->convert(d_buf, detections);
  if (detections.empty())
    std::cout << "No detection." << std::endl;
    cvtColor(h_image,h_image,COLOR_GRAY2BGR);
    for(int i = 0; i < detections.size(); ++i)
    {
      rectangle(h_image, detections[i], Scalar(0,255,255), 5);
    }

    imshow("Result image", h_image);

    waitKey(0); 
    return 0;
  }
}

代码与面部检测的代码类似。这是使用 Haar 级联的优势。如果给定对象的已训练 Haar 级联的 XML 文件可用，则相同的代码可以在所有应用中工作。只需在创建CascadeClassifier类的对象时更改 XML 文件的名称。在前面的代码中，使用了用于眼睛检测的已训练 XML 文件haarcascade_eye.xml。其他代码是自解释的。缩放因子设置为1.02，以便在每次缩放时图像大小将减少1.02。

眼睛检测程序的输出如下：

由于捕获图像时采用的角度不同，眼睛的大小各不相同，但 Haar 级联仍然能够有效地定位两只眼睛。代码的性能也可以进行测量，以查看其工作速度有多快。

总结来说，在本节中，我们展示了使用 Haar 级联进行面部和眼睛检测的应用。一旦有了训练文件，实现起来非常简单，它是一个非常强大的算法。它在内存和处理能力有限的嵌入式或移动环境中被广泛使用。

使用背景减法进行对象跟踪

背景减法是从一系列视频帧中分离前景对象与背景的过程。它在对象检测和跟踪应用中被广泛使用，以去除背景部分。背景减法分为四个步骤：

图像预处理
背景建模
前景检测
数据验证

图像预处理始终执行以去除图像中存在的任何类型的噪声。第二步是建模背景，以便它可以与前景分离。在某些应用中，视频的第一帧被用作背景，并且不进行更新。通过计算每一帧与第一帧之间的绝对差值来分离前景和背景。

在其他技术中，背景是通过取算法看到的所有帧的平均值或中值来建模的，并且该背景与前景分离。这种方法对于光照变化将更加稳健，并且会产生比第一种方法更动态的背景。还可以使用更统计密集的模型，如使用帧历史记录的高斯模型和支持向量模型来建模背景。

第三步是通过计算当前帧与背景之间的绝对差值来将前景从建模的背景中分离出来。这个绝对差值与设定的阈值进行比较，如果它大于阈值，则认为物体是移动的；如果它小于阈值，则认为物体是静止的。

高斯混合（MoG）方法

MoG 是一种广泛使用的背景减法方法，用于根据高斯混合将前景从背景中分离出来。背景从帧序列中持续更新。使用 K 个高斯分布的混合来分类像素为前景或背景。帧的时间序列也被加权以改进背景建模。持续变化的强度被分类为前景，而静态的强度被分类为背景。

OpenCV 和 CUDA 提供了一个简单的 API 来实现 MoG 背景减法。相应的代码如下：

#include <iostream>
#include <string>
#include "opencv2/opencv.hpp"
using namespace std;
using namespace cv;
using namespace cv::cuda;
int main()
{
  VideoCapture cap("abc.avi");
  if (!cap.isOpened())
  {
    cerr << "can not open camera or video file" << endl;
    return -1;
  }
  Mat frame;
  cap.read(frame);
  GpuMat d_frame;
  d_frame.upload(frame);
  Ptr<BackgroundSubtractor> mog = cuda::createBackgroundSubtractorMOG();
  GpuMat d_fgmask,d_fgimage,d_bgimage;
  Mat h_fgmask,h_fgimage,h_bgimage;
  mog->apply(d_frame, d_fgmask, 0.01);
  while(1)
  {
    cap.read(frame);
    if (frame.empty())
      break;
    d_frame.upload(frame);
    int64 start = cv::getTickCount();
    mog->apply(d_frame, d_fgmask, 0.01);
    mog->getBackgroundImage(d_bgimage);
    double fps = cv::getTickFrequency() / (cv::getTickCount() - start);
    std::cout << "FPS : " << fps << std::endl;
    d_fgimage.create(d_frame.size(), d_frame.type());
    d_fgimage.setTo(Scalar::all(0));
    d_frame.copyTo(d_fgimage, d_fgmask);
    d_fgmask.download(h_fgmask);
    d_fgimage.download(h_fgimage);
    d_bgimage.download(h_bgimage);
    imshow("image", frame);
    imshow("foreground mask", h_fgmask);
    imshow("foreground image", h_fgimage);
    imshow("mean background image", h_bgimage);
    if (waitKey(1) == 'q')
      break;
  }

  return 0;
}

使用createBackgroundSubtractorMOG类来创建用于 MoG 实现的对象。在创建对象时可以提供一些可选参数。这些参数包括history、nmixtures、backgroundRatio和noiseSigma。history参数表示用于建模背景的先前帧的数量。其默认值是 200。nmixture参数指定用于分离像素的高斯混合的数量。其默认值是 5。您可以根据应用程序的需要调整这些值。

创建的对象的apply方法用于从第一帧创建前景掩码。它需要一个输入图像和一个图像数组作为输入来存储前景掩码和学习率。在while循环的每一帧之后，都会持续更新前景掩码和背景图像。getBackgroundImage函数用于获取当前的背景模型。

前景掩码用于创建一个前景图像，指示哪些对象目前正在移动。它基本上是逻辑的，在原始帧和前景掩码之间操作。前景掩码、前景图像和建模的背景在每帧之后下载到主机内存，以便在屏幕上显示。

MoG 模型应用于 PETS 2009 数据集的视频，该数据集广泛用于行人检测。它具有静态背景，视频中有人员在移动。视频的两个不同帧的输出如下：

如所示，MoG 非常有效地建模背景。只有移动的人存在于前景掩码和前景图像中。此前景图像可用于检测到的对象的进一步处理。如果一个人停止行走，那么他将成为背景的一部分，如第二帧的结果所示。因此，此算法只能用于检测移动对象。它不会考虑静态对象。MoG 在帧率方面的性能如下：

每帧更新帧率。可以看出，它大约是每秒 330 帧，这非常高，易于用于实时应用。OpenCV 和 CUDA 还提供了 MoG 的第二版本，可以通过createBackgroundSubtractorMOG2类调用。

GMG 用于背景减法

GMG 算法的名称 GMG 来源于提出该算法的发明者的首字母。该算法是背景估计和每像素贝叶斯分割的组合。它使用贝叶斯推理将背景与前景分开。它还使用帧的历史记录来建模背景。它再次根据帧的时间序列进行加权。新的观测值比旧的观测值加权更多。

OpenCV 和 CUDA 为 GMG 算法的实现提供了与 MoG 类似的 API。实现背景减法的 GMG 算法的代码如下：

#include <iostream>
#include <string>
#include "opencv2/opencv.hpp"
#include "opencv2/core.hpp"
#include "opencv2/core/utility.hpp"
#include "opencv2/cudabgsegm.hpp"
#include "opencv2/cudalegacy.hpp"
#include "opencv2/video.hpp"
#include "opencv2/highgui.hpp"

using namespace std;
using namespace cv;
using namespace cv::cuda;

int main()
{
  VideoCapture cap("abc.avi");
  if (!cap.isOpened())
  {
    cerr << "can not open video file" << endl;
    return -1;
  }
  Mat frame;
  cap.read(frame);
  GpuMat d_frame;
  d_frame.upload(frame);
  Ptr<BackgroundSubtractor> gmg = cuda::createBackgroundSubtractorGMG(40);
  GpuMat d_fgmask,d_fgimage,d_bgimage;
  Mat h_fgmask,h_fgimage,h_bgimage;
  gmg->apply(d_frame, d_fgmask);
  while(1)
  {
    cap.read(frame);
    if (frame.empty())
      break;
    d_frame.upload(frame);
    int64 start = cv::getTickCount();
    gmg->apply(d_frame, d_fgmask, 0.01);
    double fps = cv::getTickFrequency() / (cv::getTickCount() - start);
    std::cout << "FPS : " << fps << std::endl;
    d_fgimage.create(d_frame.size(), d_frame.type());
    d_fgimage.setTo(Scalar::all(0));
    d_frame.copyTo(d_fgimage, d_fgmask);
    d_fgmask.download(h_fgmask);
    d_fgimage.download(h_fgimage);
    imshow("image", frame);
    imshow("foreground mask", h_fgmask);
    imshow("foreground image", h_fgimage);
    if (waitKey(30) == 'q')
      break;
  }
  return 0;
}

使用createBackgroundSubtractorGMG类创建用于 GMG 实现的对象。在创建对象时可以提供两个参数。第一个参数是用于建模背景的先前帧数。在上述代码中取为40。第二个参数是决策阈值，用于将像素分类为前景。其默认值为 0.8。

创建的对象的apply方法用于第一帧以创建前景掩码。通过使用帧的历史记录，前景掩码和前景图像在while循环内部持续更新。前景掩码用于以类似于 MoG 所示的方式创建前景图像。GMG 算法在相同视频和两个帧上的输出如下：

与 MoG 相比，GMG 的输出噪声更大。可以对 GMG 的结果应用形态学开闭操作，以去除结果中存在的阴影噪声。GMG 算法在 FPS 方面的性能如下：

由于它比 MoG 计算量更大，帧率较低，但仍然达到 120 FPS，这比实时性能所需的 30 FPS 要高。

总结来说，在本节中我们看到了两种背景建模和背景减法的方法。与 GMG 算法相比，MoG 算法更快且噪声更少。GMG 算法需要形态学操作来去除结果中存在的噪声。

摘要

本章描述了 OpenCV 和 CUDA 在实时目标检测和跟踪应用中的作用。它从目标检测和跟踪的介绍开始，包括在这个过程中遇到的问题和它的应用。不同的特征，如颜色、形状、直方图和其他独特的关键点，如角点，可以用来检测和跟踪图像中的对象。基于颜色的目标检测更容易实现，但要求对象与背景有明显的颜色差异。对于基于形状的目标检测，已经描述了 Canny 边缘检测技术来检测边缘，以及 Hough 变换用于直线和圆的检测。它有许多应用，如土地检测、球跟踪等。颜色和形状是全局特征，更容易计算且需要的内存较少。它们更容易受到噪声的影响。其他算法如 FAST、ORB 和 SURF 已经详细描述，这些算法可以用来从图像中检测关键点，这些关键点可以用来准确描述图像，进而可以用来检测图像中的对象。ORB 是开源的，并且无需成本就能提供与 SURF 相当的结果。SURF 是专利的，但它更快，具有尺度不变性和旋转不变性。已经描述了 Haar 级联，这是一个简单的算法，用于从图像中检测对象，如人脸、眼睛和人体。它可以用于嵌入式系统中的实时应用。本章的最后部分详细描述了背景减法算法，如 MoG 和 GMG，这些算法可以将前景与背景分离。这些算法的输出可以用于目标检测和跟踪。下一章将描述如何将这些应用部署在嵌入式开发板上。

问题

编写一个 OpenCV 代码，用于从视频中检测黄色对象。
在哪些情况下，使用颜色进行的目标检测会失败？
为什么 Canny 边缘检测算法比上一章中看到的其他边缘检测算法更好？
可以采取什么措施来降低 Hough 变换的噪声敏感性？
在 FAST 关键点检测器中，阈值的重要性是什么？
在 SURF 检测器中，Hessian 阈值的重要性是什么？
如果 Haar 级联中的尺度因子从 1.01 变为 1.05，那么它将对输出产生什么影响？
比较 MoG 和 GMG 背景减法方法。如何从 GMG 输出中去除噪声？

第八章：Jetson TX1 开发板简介及在 Jetson TX1 上安装 OpenCV

上一章介绍了使用 OpenCV 和 CUDA 的各种计算机视觉应用。当这些应用需要在实际场景中部署时，就需要一个嵌入式开发板，能够通过利用 OpenCV 和 CUDA 高速处理图像。Nvidia 提供了多个基于 GPU 的开发板，如 Jetson TK1、TX1 和 TX2，它们非常适合高端计算任务，如计算机视觉。本章将介绍其中一个开发板，即 Jetson TX1。还将详细讨论该板可用的特性和应用。CUDA 和 OpenCV 对于计算机视觉应用至关重要，因此本章将详细讨论在 Jetson TX1 上安装它们的步骤。

本章将涵盖以下主题：

Jetson TX1 开发板简介
Jetson TX1 开发板的特性和应用
基本要求和在 Jetson TX1 开发板上安装 JetPack 的步骤

技术要求

本章要求对 Linux 操作系统（OS）和网络有良好的理解。它还需要任何 Nvidia GPU 开发板，例如 Jetson TK1、TX1 或 TX2。本章中使用的 JetPack 安装文件可以从以下链接下载：developer.nvidia.com/embedded/jetpack。

Jetson TX1 简介

当高端视觉计算和计算机视觉应用需要在实际场景中部署时，就需要嵌入式开发平台，这些平台能够高效地执行计算密集型任务。例如，Raspberry Pi 可以使用 OpenCV 进行计算机视觉应用和摄像头接口功能，但对于实时应用来说非常慢。专注于 GPU 制造的 Nvidia 开发了用于计算密集型任务的模块，这些模块可用于在嵌入式平台上部署计算机视觉应用，包括 Jetson TK1、Jetson TX1 和 Jetson TX2。

Jetson TK1 是一个初步的板，包含 192 个 CUDA 核心，配备 Nvidia Kepler GPU。它是三者中最便宜的。Jetson TX1 在处理速度方面处于中等水平，拥有 256 个 CUDA 核心，采用 Maxwell 架构，运行频率为 998 MHz，同时配备 ARM CPU。Jetson TX2 在处理速度和价格方面都是最高的。它包含 256 个 CUDA 核心，采用 Pascal 架构，运行频率为 1,300 MHz。本章将详细介绍 Jetson TX1。

Jetson TX1 是专为要求较高的嵌入式应用开发的小型模块化系统。它是基于 Linux 的，提供具有万亿次浮点运算性能的超高性能，可用于计算机视觉和深度学习应用。以下照片展示了 Jetson TX1 模块：

该模块的尺寸为 50 x 87 mm，这使得它很容易集成到任何系统中。Nvidia 还提供了 Jetson TX1 开发板，该板可以快速用于原型设计应用，并包含此 GPU。以下照片展示了整个开发套件：

如照片所示，除了 GPU 模块外，开发套件还包含摄像头模块、USB 端口、以太网端口、散热器、风扇和天线。它支持包括 JetPack、Linux for Tegra、CUDA Toolkit、cuDNN、OpenCV 和 VisionWorks 在内的软件生态系统。这使得它非常适合进行深度学习和计算机视觉研究的开发者进行快速原型设计。以下章节将详细介绍 Jetson TX1 开发套件的功能。

Jetson TX1 的重要特性

Jetson TX1 开发套件具有许多特性，使其非常适合超级计算任务：

它是采用 20 nm 技术构建的系统级芯片，包含一个 1.73 GHz 的 ARM Cortex A57 四核 CPU 和一个 998 MHz 的 256 核心 Maxwell GPU。
它配备了 4 GB 的 DDR4 内存，数据总线为 64 位，工作速度为 1,600 MHz，相当于 25.6 GB/s。
它包含一个 500 万像素的 MIPI CSI-2 摄像头模块。它支持高达六路双通道或三路四通道摄像头，速度为 1,220 MP/s。
开发套件还包含一个标准 USB 3.0 类型 A 端口和微型 USB 端口，用于将鼠标、键盘和 USB 摄像头连接到板上。
它还配备了以太网端口和 Wi-Fi 连接，用于网络连接。
它可以通过 HDMI 端口连接到 HDMI 显示设备。
该套件包含一个散热器和风扇，用于在 GPU 设备达到峰值性能时进行冷却。
在空闲状态下，它仅消耗 1 瓦特的电力，在正常负载下大约为 8-10 瓦，当模块完全使用时可达 15 瓦。在 5.7 瓦的功耗下，它可以每秒处理 258 张图像，相当于性能/瓦特值为 45。一个普通的 i7 CPU 处理器在 62.5 瓦的功耗下，每秒可以处理 242 张图像，相当于性能/瓦特值为 3.88。因此，Jetson TX1 比 i7 处理器好 11.5 倍。

Jetson TX1 的应用

Jetson TX1 可用于许多需要计算密集型任务的深度学习和计算机视觉应用。以下是一些 Jetson TX1 可以使用的领域和应用：

它可用于构建各种计算密集型任务的自主机器和自动驾驶汽车。
它可用于各种计算机视觉应用，如目标检测、分类和分割。它还可以用于医学成像，分析 MRI 图像和计算机断层扫描（CT）图像。
它可以用来构建智能视频监控系统，这些系统能够帮助进行犯罪监控或交通监控。
它可以用于生物信息学和计算化学，用于模拟 DNA 基因、测序、蛋白质对接等。
它可以用于需要快速计算的各种防御设备。

在 Jetson TX1 上安装 JetPack

当 TX1 首次启动时，应该安装预装的 Linux 操作系统。以下命令可以完成安装：

cd ${HOME}/NVIDIA-INSTALLER
sudo ./installer.sh

在执行这两个命令后重启 TX1，将启动带有用户界面的 Linux 操作系统。Nvidia 提供了一套软件开发工具包（SDK），其中包含构建计算机视觉和深度学习应用所需的所有软件，以及用于闪存开发板的目标操作系统。这个 SDK 被称为JetPack。最新的 JetPack 包含 Linux for Tegra (L4T)板支持包；TensorRT，用于计算机视觉应用中的深度学习推理；最新的 CUDA 工具包，cuDNN，这是一个 CUDA 深度神经网络库；VisionWorks，它也用于计算机视觉和深度学习应用；以及 OpenCV。

当你安装 JetPack 时，所有这些包都将默认安装。本节描述了在板上安装 JetPack 的步骤。这个过程很长、繁琐，对于一个 Linux 新手来说有点复杂。所以，请仔细遵循以下章节中给出的步骤和截图。

安装的基本要求

在 TX1 上安装 JetPack 有一些基本要求。JetPack 不能直接在板上安装，因此需要一个运行 Ubuntu 14.04 的 PC 或虚拟机作为宿主机。安装过程不使用最新的 Ubuntu 版本，但你可以在它上面自由尝试。Jetson TX1 板需要一些外围设备，如鼠标、键盘和显示器，这些可以通过 USB 和 HDMI 端口连接。Jetson TX1 板应通过以太网线连接到与宿主机相同的路由器。安装还需要一根 micro USB 到 USB 线，用于通过串行传输将板与 PC 连接，以便在板上传输软件包。通过检查路由器配置来记录板的 IP 地址。如果所有要求都得到满足，则转到以下章节进行 JetPack 的安装。

安装步骤

本节描述了安装最新 JetPack 版本的步骤，并附有截图。所有步骤都需要在运行 Ubuntu 14.04 的宿主机上执行：

通过以下链接从官方 Nvidia 网站下载最新的 JetPack 版本，developer.nvidia.com/embedded/jetpack，并点击下载按钮，如图所示：

在本书撰写时，使用的最新版本是 JetPack 3.3。它用于演示安装过程。下载文件的名称是 JetPack-L4T-3.3-linux-x64_b39.run。
在桌面上创建一个名为 jetpack 的文件夹，并将此文件复制到该文件夹中，如图所示：

通过右键单击并选择“打开”选项在该文件夹中启动一个终端。该文件需要执行，因此它应该具有执行权限。如果不是这种情况，请更改权限，然后启动安装程序，如图所示：

系统将启动 JetPack 3.3 的安装向导，如图所示。只需在此窗口中点击“下一步”：

向导将询问要下载和安装软件包的目录。您可以选择当前目录进行安装，并在该目录中创建一个新文件夹以保存下载的软件包，如图所示。然后点击“下一步”：

安装向导将要求您选择要安装 JetPack 软件包的开发板。选择 Jetson TX1，如图所示，然后点击“下一步”：

组件管理器窗口将显示，显示哪些软件包将被下载和安装。它将显示 CUDA Toolkit、cuDNN、OpenCV 和 VisionWorks 等软件包，以及操作系统镜像，如图所示：

系统将要求您接受许可协议。因此，点击“接受所有”，如图所示，然后点击“下一步”：

系统将开始下载软件包，如图所示：

当所有软件包都下载并安装完成后，点击“下一步”以在主机上完成安装。它将显示以下窗口：

系统将要求您选择网络布局，即如何将板连接到主机 PC。板和主机 PC 连接到同一路由器，因此选择第一个选项，该选项指示设备通过同一路由器或交换机访问互联网，如图所示，然后点击“下一步”：

系统将询问用于将板连接到网络的接口。我们必须使用以太网线将路由器连接到板，因此我们将选择 eth0 接口，如图所示：

这将完成主机上的安装，并显示将要传输和安装到板上的包的摘要。当您在窗口中点击“下一步”时，它将显示通过微型 USB 到 USB 线将板连接到 PC 以及以强制 USB 恢复模式启动板的步骤。以下窗口显示了这些步骤：

要进入强制恢复模式，在按下电源按钮后，按下强制恢复按钮，同时按下并释放重置按钮。然后释放强制恢复按钮。设备将以强制恢复模式启动。
在窗口中输入lsusb命令；如果连接正确，它将开始将包传输到设备。如果您使用的是虚拟机，那么您必须从虚拟机的 USB 设置中启用设备。如果尚未选择，请选择 USB 3.0 控制器。在输入lsusb命令后启动的过程如下所示：

该过程将在设备上刷新操作系统。这个过程可能需要很长时间，最长可达一小时才能完成。刷新完成后，会要求重置设备以获取ssh的 IP 地址。写下之前记录的 IP 地址，以及默认的用户名和密码，即ubuntu，然后点击“下一步”。之后的窗口如下所示：

点击“下一步”，它将所有包，如 CUDA Toolkit、VisionWorks、OpenCV 和多媒体，推送到设备。以下窗口将显示：

在过程完成后，会询问是否要删除在过程中下载的所有包。如果您想删除，则勾选复选框或保持原样，如下面的截图所示：

点击“下一步”，安装过程将完成。
重启 Jetson TX1 开发板，它将以正常的 Ubuntu OS 启动。您还将观察到已安装的所有包的示例。我们将在下一章中看到如何在板上使用 CUDA 和 OpenCV。

摘要

本章介绍了用于在嵌入式平台上部署计算机视觉和深度学习应用的 Jetson TX1 开发板。它是一个小型信用卡大小的模块，可用于计算密集型应用。它的每瓦性能比最新的 i7 处理器更好。它可用于许多领域，在这些领域中，计算机视觉和深度学习被用于性能提升和嵌入式部署。Nvidia 提供了一套开发套件，其中包含此模块以及其他外围设备，可用于所有应用的快速原型设计。Nvidia 还提供了一套名为 JetPack 的 SDK，它是一系列软件包的集合，例如OpenCV、CUDA和Visionworks。本章详细描述了在 Jetson TX1 上安装 JetPack 的过程。下一章将描述使用 OpenCV 和 CUDA 在 Jetson TX1 上部署计算机视觉应用的过程。

问题

使用 Jetson TX1 而不是 Raspberry Pi 的优势是什么？
可以与 Jetson TX1 接口连接多少个摄像头？
如何连接超过两个 USB 设备到 Jetson TX1？
对或错：Jetson TX1 在功耗方面比最新的 i7 处理器有更好的性能。
对或错：Jetson TX1 不包含 CPU。
当 Jetson TX1 预装了 Ubuntu 操作系统时，JetPack 的安装要求是什么？

第九章：在 Jetson TX1 上部署计算机视觉应用

上一章介绍了在 Jetson TX1 开发板上安装 OpenCV 和 CUDA 的过程。本章将描述如何使用这些功能。将详细描述 Jetson TX1 GPU 的属性，这些属性使其适用于并行处理。本章还将描述如何在我们这本书中之前看到的 CUDA 和 C++代码在 Jetson TX1 上执行。它还将展示 Jetson TX1 GPU 在执行 CUDA 代码时的性能。本章的主要动机将是展示如何使用 Jetson TX1 部署图像和视频处理应用。以基本的图像处理应用，如图像读取、显示、加法、阈值和滤波为例，来展示如何使用 Jetson TX1 进行计算机视觉应用。此外，摄像头接口对于在实际场景中部署该板非常重要。本章将描述使用板载摄像头或 USB 摄像头进行视频捕获和处理应用的步骤。本章的最后部分将解释如何部署一些高级应用，如人脸检测和背景减法。

本章将涵盖以下主题：

Jetson TX1 板的设备属性
在 Jetson TX1 板上运行 CUDA 程序
在 Jetson TX1 板上进行图像处理
将摄像头与 Jetson TX1 开发板连接
在 Jetson TX1 开发板上执行高级应用，如人脸检测、眼检测和背景减法

技术要求

本章需要具备对 OpenCV、CUDA 和任何编程语言的良好理解。它还需要任何 Nvidia GPU 开发板，如 Jetson TK1、TX1 或 TX2。本章使用的代码文件可以从以下 GitHub 链接下载：github.com/PacktPublishing/Hands-On-GPU-Accelerated-Computer-Vision-with-OpenCV-and-CUDA。

观看以下视频以查看代码的实际运行情况：

bit.ly/2xDtHhm

Jetson TX1 GPU 的设备属性

CUDA 提供了一个简单的接口来确定 GPU 设备的性能，这是 Jetson TX1 板上存在的 Tegra X1。了解设备的属性对于编写针对它的优化程序非常重要。查找设备属性的程序包含在 JetPack 安装的 CUDA 示例程序中，位于主文件夹中。您还可以运行我们在第二章中开发的程序来查找设备属性。

程序在 Nvidia Tegra X1 GPU 上的输出如下：

JetPack 3.3 安装了 CUDA 9.0 运行时版本。GPU 设备的全局内存大约为 4 GB，GPU 时钟速度约为 1 GHz。这个时钟速度比本书前面提到的 GeForce 940 GPU 慢。内存时钟速度仅为 13 MHz，而 GeForce 940 为 2.505 GHz，这使得 Jetson TX1 较慢。与 GeForce 940 的 1 MB 相比，L2 缓存为 256 KB。大多数其他属性与 GeForce 940 相似。

在 X、Y 和 Z 方向上，每个块可以启动的最大线程数分别为 1,024、1,024 和 64。在确定从程序中启动的并行线程数量时，应使用这些数字。在启动每个网格的并行块数量时，也应采取相同的谨慎措施。

总结来说，我们已经看到了 Jetson TX1 开发板上可用的 Tegra X1 GPU 的设备属性。它是一个嵌入式板，因此内存可用，时钟速度相对于笔记本电脑中配备的 GPU 设备（如 GeForce 940）较慢。然而，它比 Arduino 和 Raspberry Pi 等嵌入式平台快得多。它可以很容易地用于部署需要高性能计算能力的计算机视觉应用。现在我们已经看到了设备属性，我们将从在 Jetson TX1 上使用 CUDA 开发第一个程序开始。

基于 Jetson TX1 的 CUDA 基础程序

在本节中，我们通过添加两个大数组的示例来展示使用 Jetson TX1 开发板执行 CUDA 程序。程序的性能也使用 CUDA 事件进行了测量。

添加具有 50,000 个元素的两个大数组的内核函数如下：


#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
//Defining number of elements in Array
#define N 50000
//Defining Kernel function for vector addition
__global__ void gpuAdd(int *d_a, int *d_b, int *d_c) {
 //Getting Thread index of current kernel
 int tid = threadIdx.x + blockIdx.x * blockDim.x;
 while (tid < N)
 {
 d_c[tid] = d_a[tid] + d_b[tid];
 tid += blockDim.x * gridDim.x;
 }
}

内核函数接收两个设备指针，这些指针指向输入数组作为输入，以及一个设备指针，该指针指向设备内存中输出数组的参数。计算当前内核执行的线程 ID，并由内核将线程索引索引的数组元素相加。如果启动的内核数量少于数组元素数量，则相同的内核将在while循环中添加由块维度偏移的Array元素。添加两个数组的main函数如下：

int main(void) 
{
 //Defining host arrays
 int h_a[N], h_b[N], h_c[N];
 //Defining device pointers
 int *d_a, *d_b, *d_c;
 cudaEvent_t e_start, e_stop;
 cudaEventCreate(&e_start);
 cudaEventCreate(&e_stop);
 cudaEventRecord(e_start, 0);
 // allocate the memory
 cudaMalloc((void**)&d_a, N * sizeof(int));
 cudaMalloc((void**)&d_b, N * sizeof(int));
 cudaMalloc((void**)&d_c, N * sizeof(int));
 //Initializing Arrays
 for (int i = 0; i < N; i++) {
 h_a[i] = 2 * i*i;
 h_b[i] = i;
 }
 // Copy input arrays from host to device memory
 cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice);
 cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice);
 //Calling kernels passing device pointers as parameters
 gpuAdd << <1024, 1024 >> >(d_a, d_b, d_c);
 //Copy result back to host memory from device memory
 cudaMemcpy(h_c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);
 cudaDeviceSynchronize();
 cudaEventRecord(e_stop, 0);
 cudaEventSynchronize(e_stop);
 float elapsedTime;
 cudaEventElapsedTime(&elapsedTime, e_start, e_stop);
 printf("Time to add %d numbers: %3.1f ms\n",N, elapsedTime);

定义了两个主机数组，并使用cudaMalloc函数为它们分配内存。它们被初始化为一些随机值，并上传到设备内存。创建了两个 CUDA 事件来测量 CUDA 程序的性能。内核以并行方式启动了 1,024 个块，每个块有 1,024 个线程。这些数字来自设备属性，如上一节所述。内核函数的结果被传输到主机内存。内核函数的时间由e_start和e_stop事件在内核启动前后记录。函数所花费的时间显示在控制台上。

以下代码被添加以验证由 GPU 计算的结果的正确性，并清理程序使用的内存：

 int Correct = 1;
 printf("Vector addition on GPU \n");
 //Printing result on console
 for (int i = 0; i < N; i++) {
 if ((h_a[i] + h_b[i] != h_c[i]))
 {
  Correct = 0;
 }

 }
 if (Correct == 1)
 {
 printf("GPU has computed Sum Correctly\n");
 }
 else
 {
 printf("There is an Error in GPU Computation\n");
 }
 //Free up memory
 cudaFree(d_a);
 cudaFree(d_b);
 cudaFree(d_c);
 return 0;
}

在 CPU 上执行相同的数组加法操作，并与从 GPU 获得的结果进行比较，以验证 GPU 是否正确计算了结果。这也在控制台上显示。通过使用cudaFree函数释放程序使用的所有内存。

需要从终端运行以下两个命令来执行程序。程序应位于当前工作目录中：

$ nvcc 01_performance_cuda_events.cu -o gpu_add
$ ./gpu_add

使用nvcc命令通过 Nvidia CUDA 编译器编译 CUDA 代码。文件名作为命令的参数传递。编译器将创建的目标文件名通过-o选项指定。此文件名将用于执行程序。这是通过第二个命令完成的。程序输出如下：

从结果可以看出，Jetson TX1 计算包含 50,000 个元素的两个数组的和需要3.4ms，这比本书第三章中使用的 GeForce 940 慢，但仍然比 CPU 上的顺序执行快。

总结来说，本节展示了在执行 CUDA 程序中使用 Jetson TX1 开发板的方法。语法与我们在这本书中之前看到的相同。因此，书中之前开发的全部 CUDA 程序都可以在 Jetson TX1 上执行，无需太多修改。程序执行的步骤也进行了描述。下一节将描述使用 Jetson TX1 进行图像处理应用的方法。

在 Jetson TX1 上进行的图像处理

本节将展示在部署图像处理应用中使用 Jetson TX1 的方法。我们还将再次使用 OpenCV 和 CUDA 来加速 Jetson TX1 上的计算机视觉应用。在上一章中，我们看到了包含 OpenCV 和 CUDA 的 JetPack 3.3 的安装过程。但在最新的 JetPack 中，OpenCV 没有编译带有 CUDA 支持，也没有 GStreamer 支持，这是从代码中访问摄像头所必需的。因此，删除 JetPack 附带的 OpenCV 安装，并编译带有 CUDA 和 GStreamer 支持的 OpenCV 新版本是一个好主意。下一节将展示如何执行此过程。

（编译带有 CUDA 支持的 OpenCV，如果需要）

尽管 JetPack 附带的 OpenCV 可以与新安装的 OpenCV 一起工作，但先删除旧安装然后开始新的安装是一个好主意。这将避免不必要的混淆。为了完成这个任务，必须执行以下步骤：

从终端运行以下命令：

$ sudo apt-get purge libopencv*

确保安装的所有软件包都是最新版本。如果不是这样，可以通过运行以下两个命令来更新它们：

$ sudo apt-get update
$ sudo apt-get dist-upgrade

编译从源代码编译 OpenCV 需要最新的 cmake 和 gcc 编译器版本，因此可以通过运行以下两个命令来安装：

$ sudo apt-get install --only-upgrade gcc-5 cpp-5 g++-5
$ sudo apt-get install build-essential make cmake cmake-curses-gui libglew-dev libgtk2.0-dev

需要安装一些依赖项以编译支持 GStreamer 的 OpenCV。这可以通过以下命令完成：

sudo apt-get install libdc1394-22-dev libxine2-dev libgstreamer1.0-dev libgstreamer-plugins-base1.0-dev

通过执行以下命令下载 OpenCV 最新版本的源代码，并将其提取到一个文件夹中：

$ wget https://github.com/opencv/opencv/archive/3.4.0.zip -O opencv.zip
$ unzip opencv.zip

现在，进入 opencv 文件夹，创建 build 目录。然后进入这个新创建的 build 目录。这些可以通过从命令提示符执行以下命令来完成。

$ cd opencv
$ mkdir build
$ cd build

cmake 命令用于编译具有 CUDA 支持的 opencv。确保在此命令中将 WITH_CUDA 标志设置为 ON。注意，对于 Jetson TX1 开发板，CUDA_ARCH_BIN 应设置为 5.3，对于 Jetson TX2，应设置为 6.2。示例没有构建以节省时间和空间。整个 cmake 命令如下：

cmake -D CMAKE_BUILD_TYPE=RELEASE -D CMAKE_INSTALL_PREFIX=/usr/local \
 -D WITH_CUDA=ON -D CUDA_ARCH_BIN="5.3" -D CUDA_ARCH_PTX="" \
 -D WITH_CUBLAS=ON -D ENABLE_FAST_MATH=ON -D CUDA_FAST_MATH=ON \
 -D ENABLE_NEON=ON -D WITH_LIBV4L=ON -D BUILD_TESTS=OFF \
 -D BUILD_PERF_TESTS=OFF -D BUILD_EXAMPLES=OFF \
 -D WITH_QT=ON -D WITH_OPENGL=ON ..

它将启动配置和创建 makefile 的过程。在配置成功后，cmake 命令将在 build 目录中创建 makefile。
要使用 makefile 编译 OpenCV，请在命令窗口中执行 make -j4 命令。
编译成功后，您必须从命令行执行 sudo make install 命令来安装 OpenCV。

如果这些步骤执行成功，则 OpenCV 3.4.0 将在 Jetson TX1 上安装，并支持 CUDA 和 GStreamer，使用 OpenCV 制作的任何计算机视觉应用都可以部署到它上面。下一节将演示在板上进行简单的图像处理操作。

读取和显示图像

对于任何计算机视觉应用，所需的基本操作之一是读取和显示存储在磁盘上的图像。本节将演示一个简单的代码，在 Jetson TX1 上执行此操作。当我们从计算机上的 GPU 移动到 Jetson TX1 开发板时，OpenCV 语法不会改变太多。将有一些小的变化。在 Jetson TX1 上读取和显示图像的代码如下：

#include <opencv2/opencv.hpp>
#include <iostream>

using namespace cv;
using namespace std;

int main()
{
 Mat img = imread("images/cameraman.tif",0);
 if (img.empty()) 
 {
 cout << "Could not open an image" << endl;
 return -1;
 }
 imshow("Image Read on Jetson TX1"; , img); 
 waitKey(0); 
 return 0;
}

必要的 OpenCV 库包含在代码中。图像是通过 Main 函数内的 imread 函数读取的。由于 imread 命令的第二个参数指定为 0，因此图像被读取为灰度图像。要将图像读取为彩色图像，可以指定为 1。if 语句检查图像是否被读取，如果没有，则在控制台上显示错误后终止代码。当图像名称不正确或图像未存储在指定的路径时，可能会发生读取图像的错误。这个错误由 if 语句处理。图像是通过 imshow 命令显示的。waitKey 函数用于显示图像，直到按下键盘上的任何键。

所示的前面代码可以保存为image_read.cpp文件，并使用以下命令从终端执行。请确保程序文件存储在终端的当前工作目录中：

For compilation:
$ g++ -std = c++11 image_read.cpp 'pkg_config --libs --cflags opencv' -o image_read
For execution:
$./image_read

程序的输出如下：

本节演示了在 Jetson TX1 上读取和显示图像的步骤。在下节中，我们将看到一些更多的图像处理操作，并尝试测量它们在 Jetson TX1 上的性能。

图像加法

本节将演示使用 Jetson TX1 进行简单的图像处理应用，如图像加法。在加法后，相同位置的像素强度被相加以构建新的图像。假设在两个图像中，(0,0)位置的像素强度值分别为 50 和 150，那么结果图像中的强度值将是 200，这是两个强度值的和。OpenCV 的加法操作是饱和操作，这意味着如果加法的结果超过 255，则将饱和在 255。在 Jetson TX1 上执行加法的代码如下：

#include <iostream>
#include "opencv2/opencv.hpp"
#include "opencv2/core/cuda.hpp"

int main (int argc, char* argv[])
{
 //Read Two Images 
 cv::Mat h_img1 = cv::imread("images/cameraman.tif");
 cv::Mat h_img2 = cv::imread("images/circles.png");
 int64 work_begin = cv::getTickCount(); 
 //Create Memory for storing Images on device
 cv::cuda::GpuMat d_result1,d_img1, d_img2;
 cv::Mat h_result1;
 //Upload Images to device 
 d_img1.upload(h_img1);
 d_img2.upload(h_img2);

 cv::cuda::add(d_img1,d_img2, d_result1);
 //Download Result back to host
 d_result1.download(h_result1);
 cv::imshow("Image1 ", h_img1);
 cv::imshow("Image2 ", h_img2);
 cv::imshow("Result addition ", h_result1);
 int64 delta = cv::getTickCount() - work_begin;
 //Frequency of timer
 double freq = cv::getTickFrequency();
 double work_fps = freq / delta;
 std::cout<<"Performance of Addition on Jetson TX1: " <<std::endl;
 std::cout <<"Time: " << (1/work_fps) <<std::endl;
 std::cout <<"FPS: " <<work_fps <<std::endl;

 cv::imshow("result_add.png", h_result1);
 cv::waitKey();
 return 0;
}

在进行图像加法时需要注意的一点是，两个图像应该具有相同的大小。如果不是这种情况，那么在加法之前应该将它们调整大小。在前面代码中，从磁盘读取了两个相同大小的图像，并将它们上传到设备内存中进行 GPU 上的加法。使用cv::cuda模块中的add函数在设备上执行图像加法。结果图像被下载到主机并在控制台上显示。

程序的输出如下：

使用cv::getTickCount()和cv::getTickFrequency()函数也测量了图像加法的性能。加法操作所需的时间显示在下面的屏幕截图上：

如前述屏幕截图所示，在 Jetson TX1 上添加两个大小为 256 x 256 的图像大约需要0.26ms。这对于嵌入式平台来说是一个非常不错的性能。需要注意的是，在测量加法操作的准确时间之前，应该测量imshow函数。imshow函数显示图像需要更多的时间，因此测量的时间不会是加法操作所需时间的准确估计。

图像阈值化

本节将演示使用 Jetson TX1 进行更复杂的计算机视觉应用，如图像阈值化。图像阈值化是一种非常简单的图像分割技术，用于根据某些强度值从灰度图像中提取重要区域。在这种技术中，如果像素值大于某个阈值值，则分配一个值，否则分配另一个值。

OpenCV 提供了不同类型的阈值技术，这由函数的最后一个参数决定。这些阈值类型包括：

cv:.THRES H_BINARY: 如果像素的强度大于阈值，则将像素强度设置为等于maxVal常量，否则将像素强度设置为等于零。
cv::THRESH_BINARY_INV: 如果像素的强度大于阈值，则将像素强度设置为等于零，否则将像素强度设置为maxVal常量。
cv::THRESH_TRUNC: 这基本上是一个截断操作。如果像素的强度大于阈值，则将像素强度设置为等于阈值，否则保持强度值不变。
cv::THRESH_TOZERO: 如果像素的强度大于阈值，则保持像素强度不变，否则将像素强度设置为等于零。
cv::THRESH_TOZERO_INV: 如果像素的强度大于阈值，则将该像素强度设置为等于零，否则保持像素强度不变。

在 Jetson TX1 上使用 OpenCV 和 CUDA 实现所有这些阈值技术的程序如下所示：

#include <iostream>
#include "opencv2/opencv.hpp"
using namespace cv;
int main (int argc, char* argv[])
{
 cv::Mat h_img1 = cv::imread("images/cameraman.tif", 0);
 cv::cuda::GpuMat d_result1,d_result2,d_result3,d_result4,d_result5, d_img1;
 //Measure initial time ticks
 int64 work_begin = getTickCount(); 
 d_img1.upload(h_img1);
 cv::cuda::threshold(d_img1, d_result1, 128.0, 255.0, cv::THRESH_BINARY);
 cv::cuda::threshold(d_img1, d_result2, 128.0, 255.0, cv::THRESH_BINARY_INV);
 cv::cuda::threshold(d_img1, d_result3, 128.0, 255.0, cv::THRESH_TRUNC);
 cv::cuda::threshold(d_img1, d_result4, 128.0, 255.0, cv::THRESH_TOZERO);
 cv::cuda::threshold(d_img1, d_result5, 128.0, 255.0, cv::THRESH_TOZERO_INV);

 cv::Mat h_result1,h_result2,h_result3,h_result4,h_result5;
 d_result1.download(h_result1);
 d_result2.download(h_result2);
 d_result3.download(h_result3);
 d_result4.download(h_result4);
 d_result5.download(h_result5);
 //Measure difference in time ticks
 int64 delta = getTickCount() - work_begin;
 double freq = getTickFrequency();
 //Measure frames per second
 double work_fps = freq / delta;
 std::cout <<"Performance of Thresholding on GPU: " <<std::endl;
 std::cout <<"Time: " << (1/work_fps) <<std::endl;
 std::cout <<"FPS: " <<work_fps <<std::endl;
 return 0;
}

在 OpenCV 和 CUDA 的 GPU 上用于图像阈值化的函数是cv::cuda::threshold。此函数有许多参数。第一个参数是源图像，它应该是一个灰度图像。第二个参数是结果要存储的目标位置。第三个参数是阈值值，用于分割像素值。第四个参数是maxVal常量，表示如果像素值超过阈值值时赋予的值。最后一个参数是前面讨论过的阈值方法。以下程序显示了原始图像和五种阈值技术输出的输出：

使用cv::getTickCount()和cv::getTickFrequency()函数测量图像阈值化的性能。五个阈值操作所需的时间显示在控制台上，如下面的截图所示：

在 Jetson TX1 上执行五个阈值操作需要0.32ms，这对于嵌入式平台上的图像分割任务来说，性能非常好。下一节将描述 Jetson TX1 上的滤波操作。

Jetson TX1 上的图像滤波

图像滤波是图像预处理和特征提取中的一个非常重要的步骤。低通滤波器，如平均、高斯和中值滤波器，用于去除图像中的不同类型的噪声，而高通滤波器，如 Sobel、Scharr 和 Laplacian，用于检测图像中的边缘。边缘是重要的特征，可用于计算机视觉任务，如目标检测和分类。本书中已详细解释了图像滤波。

本节描述了在 Jetson TX1 上对图像应用低通和高通滤波器的步骤。相应的代码如下：

#include <iostream>
#include <string>
#include "opencv2/opencv.hpp"

using namespace std;
using namespace cv;
using namespace cv::cuda;

int main()
{
 Mat h_img1;
 cv::cuda::GpuMat d_img1,d_blur,d_result3x3;
 h_img1 = imread("images/blobs.png",1);

 int64 start = cv::getTickCount();
 d_img1.upload(h_img1);
 cv::cuda::cvtColor(d_img1,d_img1,cv::COLOR_BGR2GRAY);
 cv::Ptr<cv::cuda::Filter> filter3x3;
 filter3x3 = cv::cuda::createGaussianFilter(CV_8UC1,CV_8UC1,cv::Size(3,3),1);
 filter3x3->apply(d_img1, d_blur);

 cv::Ptr<cv::cuda::Filter> filter1;
 filter1 = cv::cuda::createLaplacianFilter(CV_8UC1,CV_8UC1,1);
 filter1->apply(d_blur, d_result3x3);

 cv::Mat h_result3x3,h_blur;
 d_result3x3.download(h_result3x3);
 d_blur.download(h_blur);

 double fps = cv::getTickFrequency() / (cv::getTickCount() - start);
 std::cout << "FPS : " << fps << std::endl;
 imshow("Laplacian", h_result3x3);
 imshow("Blurred", h_blur);
 cv::waitKey();
 return 0;
}

拉普拉斯算子是一种二阶导数，用于从图像中提取垂直和水平图像。它对噪声非常敏感，因此有时需要使用低通滤波器（如高斯模糊）来去除噪声，然后再应用拉普拉斯滤波器。因此，在代码中，使用标准差等于1的 3x3 高斯滤波器对输入图像进行处理。该滤波器是通过 OpenCV 的cv::cuda::createGaussianFilter函数创建的。然后，将拉普拉斯滤波器应用于高斯模糊后的图像。拉普拉斯滤波器是通过 OpenCV 的cv::cuda::createLaplacianFilter函数创建的。高斯模糊和拉普拉斯滤波器的输出被下载回主机内存，以便在控制台上显示。代码中还测量了滤波操作的性能。程序的输出如下截图所示：

从输出中可以看出，对模糊图像应用拉普拉斯滤波器将去除图像中的虚假边缘。它还将去除输入图像中存在的高斯噪声。如果输入图像被盐和胡椒噪声扭曲，则应在拉普拉斯滤波器进行边缘检测之前使用中值滤波器作为预处理步骤。

总结来说，我们已经在 Jetson TX1 上看到了不同的图像处理函数，如图像加法、图像阈值和图像滤波。我们还看到，这些操作在 Jetson TX1 上的性能比在 CPU 上执行相同的代码要好得多。下一节将描述如何将摄像头与 Jetson TX1 连接，以便在现实生活中的场景中使用。

与 Jetson TX1 连接摄像头

Jetson TX1 可以与 USB 摄像头或 CSI 摄像头连接。开发板已预装了一款 5 兆像素的摄像头，并与 Jetson TX1 连接。这款摄像头可以像笔记本电脑上的网络摄像头一样捕获视频。摄像头连接是 Jetson TX1 开发板在实时应用中的重要特性，它支持多达六通道的摄像头。Jetson TX1 支持的摄像头详细列表可以在以下链接中找到：elinux.org/Jetson_TX1

本节将演示使用与 Jetson TX1 连接的摄像头捕获视频的步骤，以及如何使用这些视频开发计算机视觉应用，如人脸检测和背景减法。

从机载摄像头读取和显示视频

本节将描述从 USB 摄像头或与 Jetson TX1 连接的机载摄像头捕获视频的方法。为此，OpenCV 应编译为支持 GStreamer；否则，OpenCV 将不支持捕获视频的格式。

以下代码可以用来从摄像头捕获视频并在屏幕上显示：

#include <opencv2/opencv.hpp>
#include <iostream>
#include <stdio.h>
using namespace cv;
using namespace std;

int main(int, char**)
{
 Mat frame;
 // open the default camera using default API
 VideoCapture cap("nvcamerasrc ! video/x-raw(memory:NVMM), width=(int)1280, height=(int)720, format=(string)I420, framerate=(fraction)24/1 ! nvvidconv flip-method=0 ! video/x-raw, format=(string)I420 ! videoconvert ! video/x-raw, format=(string)BGR ! appsink"); 
 if (!cap.isOpened()) {
 cout << "Unable to open camera\n";
 return -1;
 }
 while (1)
 {
 int64 start = cv::getTickCount();
 cap.read(frame);
 // check if we succeeded
 if (frame.empty()) {
  cout << "Can not read frame\n";
  break;
 }
 double fps = cv::getTickFrequency() / (cv::getTickCount() - start);
 std::cout << "FPS : " << fps << std::endl;

 imshow("Live", frame);
 if (waitKey(30) == 'q')
  break;
 }

 return 0;
}

代码与用于在台式机上从网络摄像头捕获视频的代码大致相似。而不是使用设备 ID 作为参数来捕获对象，使用指定 GStreamer 管道的字符串。如下所示：

VideoCapture cap("nvcamerasrc ! video/x-raw(memory:NVMM), width=(int)1280, height=(int)720, format=(string)I420, framerate=(fraction)24/1 ! nvvidconv flip-method=0 ! video/x-raw, format=(string)I420 ! videoconvert ! video/x-raw, format=(string)BGR ! appsink");

捕获的视频的宽度和高度被指定为 1,280 和 720 像素。帧率也被指定。这些值将根据接口摄像头的支持的格式而变化。使用nvvidconv将视频转换为 OpenCV 支持的 BGR 格式。它还用于图像缩放和翻转。要翻转捕获的视频，可以将 flip 方法指定为一个非零的整数值。

使用cap.isOpened属性来检查是否已从摄像头开始捕获。然后使用 read 方法逐个读取帧并在屏幕上显示，直到用户按下q键。代码中也测量了帧捕获的速率。

摄像头捕获了实时视频的两个不同帧，帧率显示在下面的屏幕截图上：

总结来说，在本节中，我们看到了从与 Jetson TX1 开发板相连的摄像头捕获视频的步骤。这个捕获的视频可以用于开发下一节所述的有用的实时计算机视觉应用。

Jetson TX1 上的高级应用

本节将描述在部署高级计算机视觉应用（如人脸检测、眼检测和背景减法）中使用 Jetson TX1 嵌入式平台。

使用 Haar 级联进行人脸检测

Haar 级联使用矩形特征来检测对象。它使用不同大小的矩形来计算不同的线和边缘特征。Haar-like 特征检测算法背后的思想是计算矩形内白色像素总和与黑色像素总和之间的差异。

这种方法的主要优势是使用积分图方法进行快速的求和计算。这使得 Haar 级联非常适合实时目标检测。它处理图像所需的时间比其他用于目标检测的算法要少。由于 Haar 级联具有低计算复杂性和低内存占用，因此非常适合部署在嵌入式系统如 Jetson TX1 上。因此，在本节中，使用此算法在 Jetson TX1 上部署人脸检测应用。

从与 Jetson TX1 接口的摄像头捕获的视频进行人脸检测的代码如下：

#include <iostream>
#include <opencv2/opencv.hpp>
using namespace cv;
using namespace std;

int main()
{
 VideoCapture cap("images/output.avi");
//cv::VideoCapture cap("nvcamerasrc ! video/x-raw(memory:NVMM), width=(int)1280, height=(int)720, format=(string)I420, framerate=(fraction)24/1 ! nvvidconv flip-method=0 ! video/x-raw, format=(string)I420 ! videoconvert ! video/x-raw, format=(string)BGR ! appsink"); 
 if (!cap.isOpened()) {
   cout << "Can not open video source";
   return -1;
 }
 std::vector<cv::Rect> h_found;
 cv::Ptr<cv::cuda::CascadeClassifier> cascade = cv::cuda::CascadeClassifier::create("haarcascade_frontalface_alt2.xml");
 cv::cuda::GpuMat d_frame, d_gray, d_found;
 while(1)
 {
 Mat frame;
 if ( !cap.read(frame) ) {
   cout << "Can not read frame from webcam";
   return -1;
 }
 int64 start = cv::getTickCount();
 d_frame.upload(frame);
 cv::cuda::cvtColor(d_frame, d_gray, cv::COLOR_BGR2GRAY);

 cascade->detectMultiScale(d_gray, d_found);
 cascade->convert(d_found, h_found);

 for(int i = 0; i < h_found.size(); ++i)
 {
   rectangle(frame, h_found[i], Scalar(0,255,255), 5);
 }
 double fps = cv::getTickFrequency() / (cv::getTickCount() - start);
 std::cout << "FPS : " << fps << std::endl;
 imshow("Result", frame);
 if (waitKey(1) == 'q') {
   break;
 }
 }

 return 0;
}

Haar 级联是一种需要训练以执行特定任务的算法。从头开始训练特定应用的 Haar 级联是困难的，因此 OpenCV 提供了一些训练好的 XML 文件，可用于检测对象。这些 XML 文件位于 OpenCV 和 CUDA 安装的\usr\local\opencv\data\haarcascades_cuda目录中。

初始化网络摄像头，并逐个捕获网络摄像头的帧。将帧上传到设备内存以在 GPU 上处理。OpenCV 和 CUDA 提供了CascadeClassifier类，可用于实现 Haar 级联。使用create方法创建该类的对象。它需要加载训练好的 XML 文件的文件名。

在while循环内部，将detectMultiscale方法应用于每一帧，以便在每一帧中检测不同大小的人脸。使用convert方法将检测到的位置转换为矩形向量。然后，使用for循环迭代此向量，以便在所有检测到的人脸上使用矩形函数绘制边界框。此过程对从网络摄像头捕获的每一帧重复进行。算法的性能也以每秒帧数来衡量。

程序的输出如下：

从输出中可以看出，人脸在两个不同位置的两个不同的网络摄像头帧中被正确定位。第二帧有点模糊，但这不会影响算法。Jetson TX1 上算法的性能也在右图显示。算法在大约每秒五帧的速度下工作。

总结来说，本节展示了使用 Jetson TX1 从网络摄像头捕获的实时视频中检测人脸。此应用程序可用于人员识别、人脸锁定、考勤监控等。

使用 Haar 级联进行眼检测

本节将描述使用 Haar 级联检测人类眼睛的使用。用于眼检测的训练好的 Haar 级联的 XML 文件位于 OpenCV 安装目录中。此文件用于检测眼睛。其代码如下：

#include "opencv2/objdetect/objdetect.hpp"
#include "opencv2/highgui/highgui.hpp"
#include "opencv2/imgproc/imgproc.hpp"
#include "opencv2/cudaobjdetect.hpp" 
#include <iostream>
#include <stdio.h>

using namespace std;
using namespace cv;

int main( )
{
  Mat h_image;
  h_image = imread("images/lena_color_512.tif", 0); 
  Ptr<cuda::CascadeClassifier> cascade =       cuda::CascadeClassifier::create("haarcascade_eye.xml");
  cuda::GpuMat d_image;
  cuda::GpuMat d_buf;
  int64 start = cv::getTickCount();
  d_image.upload(h_image);
  cascadeGPU->setMinNeighbors(0);
  cascadeGPU->setScaleFactor(1.02);
  cascade->detectMultiScale(d_image, d_buf);
  std::vector<Rect> detections;
  cascade->convert(d_buf, detections);
 if (detections.empty())
   std::cout << "No detection." << std::endl;
   cvtColor(h_image,h_image,COLOR_GRAY2BGR);
 for(int i = 0; i < detections.size(); ++i)
 {
   rectangle(h_image, detections[i], Scalar(0,255,255), 5);
 }
 double fps = cv::getTickFrequency() / (cv::getTickCount() - start);
 std::cout << "FPS : " << fps << std::endl;
 imshow("Result image on Jetson TX1", h_image);

 waitKey(0); 
 return 0;
}

代码与面部检测的代码类似。这是使用 Haar 级联的优势。如果有一个给定对象的训练好的 Haar 级联的 XML 文件可用，那么相同的代码将在所有应用程序中工作。只需在创建CascadeClassifier类的对象时更改 XML 文件的名称。在前面的代码中，使用了用于眼检测的训练 XML 文件haarcascade_eye.xml。其他代码是自解释的。缩放因子设置为1.02，以便在每次缩放时图像大小将减少1.02。眼检测程序的输出如下：

现在我们已经使用 Haar 级联从视频和图像中检测到对象，因此捕获的视频也可以使用下一节中描述的背景减法方法来检测和跟踪对象。

使用高斯混合（MoG）的背景减法

背景减法是目标检测和跟踪应用的重要预处理步骤。它也可以用于从监控录像中检测异常活动。本节展示了在背景减法应用中使用 Jetson TX1。与 Jetson TX1 接口的摄像头被安装在一个房间内，用于检测房间内的活动。房间的背景在第一帧中被初始化。

MoG 是一种广泛使用的背景减法方法，用于根据高斯混合将前景与背景分离，用于活动检测。背景从帧序列中持续更新。使用 K 个高斯分布的混合来将像素分类为前景或背景。帧的时间序列也被加权以改进背景建模。持续变化的强度被分类为前景，而静态的强度被分类为背景。

使用 MoG 进行活动监控的代码如下：

#include <iostream>
#include <string>
#include "opencv2/opencv.hpp"

using namespace std;
using namespace cv;
using namespace cv::cuda;
int main()
{

 VideoCapture cap("nvcamerasrc ! video/x-raw(memory:NVMM), width=(int)1280, height=(int)720, format=(string)I420, framerate=(fraction)24/1 ! nvvidconv flip-method=0 ! video/x-raw, format=(string)I420 ! videoconvert ! video/x-raw, format=(string)BGR ! appsink");
 if (!cap.isOpened())
 {
 cout << "Can not open camera or video file" << endl;
 return -1;
 }
 Mat frame;
 cap.read(frame);
 GpuMat d_frame;
 d_frame.upload(frame);
 Ptr<BackgroundSubtractor> mog = cuda::createBackgroundSubtractorMOG();
 GpuMat d_fgmask,d_fgimage,d_bgimage;
 Mat h_fgmask,h_fgimage,h_bgimage;
 mog->apply(d_frame, d_fgmask, 0.01);
 namedWindow("image", WINDOW_NORMAL);
 namedWindow("foreground mask", WINDOW_NORMAL);
 namedWindow("foreground image", WINDOW_NORMAL);
 namedWindow("mean background image", WINDOW_NORMAL);

 while(1)
 {
 cap.read(frame);
 if (frame.empty())
  break;
 d_frame.upload(frame);
 int64 start = cv::getTickCount();
 mog->apply(d_frame, d_fgmask, 0.01);
 mog->getBackgroundImage(d_bgimage);
 double fps = cv::getTickFrequency() / (cv::getTickCount() - start);
 std::cout << "FPS : " << fps << std::endl;
 d_fgimage.create(d_frame.size(), d_frame.type());
 d_fgimage.setTo(Scalar::all(0));
 d_frame.copyTo(d_fgimage, d_fgmask);
 d_fgmask.download(h_fgmask);
 d_fgimage.download(h_fgimage);
 d_bgimage.download(h_bgimage);
 imshow("image", frame);
 imshow("foreground mask", h_fgmask);
 imshow("foreground image", h_fgimage);
 imshow("mean background image", h_bgimage);
 if (waitKey(1) == 'q')
  break;
 }

 return 0;
}

与 Jetson TX1 接口的摄像头使用 GStreamer 管道初始化。createBackgroundSubtractorMOG类用于创建 MoG 实现的对象。创建的对象的apply方法用于从第一帧创建前景掩码。它需要一个输入图像、一个image数组来存储前景掩码，以及学习率作为输入。没有活动的房间图像被初始化为 MoG 的背景。因此，任何发生的活动都将被算法归类为前景。

在while循环的每一帧之后，都会持续更新前景掩码和背景图像。getBackgroundImage函数用于获取当前的背景模型。

前景掩码用于创建前景图像，指示哪些对象当前正在移动。它基本上是逻辑操作，在原始帧和前景掩码之间进行。在每一帧之后，前景掩码、前景图像和建模的背景都会下载到主机内存中，以便在屏幕上显示。

以下截图显示了从视频中提取的两个不同帧的输出：

第一行表示房间内没有任何活动时的背景。当有人将手放在摄像头前时，它将被检测为前景，如第二帧结果所示。同样，如果有人将手机放在摄像头前，它也将被归类为前景，如第三帧所示。以下截图显示了代码在每秒帧数方面的性能：

该技术在每秒大约 60-70 帧的速度下工作，可以很容易地用于实时决策。尽管本节中的演示非常简单，但这种应用可以在许多实际情况下使用。房间内的活动可以用来控制房间内现有的设备。这有助于在没有人的情况下节省电力。此应用还可以用于 ATM 机内活动的监控。它还可以用于公共场所的其他视频监控应用。Python 也可以用作 Jetson TX1 上的编程语言，这将在下一节中解释。

在 Jetson TX1 上使用 Python 和 OpenCV 进行计算机视觉

到目前为止，我们使用 C/C++、OpenCV 和 CUDA 开发了所有计算机视觉应用。Jetson TX1 也支持 Python 编程语言用于计算机视觉应用。当在 Jetson TX1 上编译 OpenCV 时，它也会安装 OpenCV 的 Python 二进制文件。因此，熟悉 Python 编程语言的程序员可以使用 Python 接口开发 OpenCV 计算机视觉应用，并在 Jetson TX1 上部署它们。Python 也像所有 Linux 操作系统一样预安装在 Jetson TX1 上。Windows 用户可以单独安装 Python。Python 的安装过程和优势将在下一章中解释。

使用 Python 的一个缺点是，OpenCV Python 接口尚未从 CUDA 加速中受益很大。尽管如此，Python 学习的简便性和其广泛的应用范围已经鼓励了许多软件开发者使用 Python 进行计算机视觉应用。使用 Python 和 OpenCV 读取和显示图像的示例代码如下：

import numpy as np
import cv2
img = cv2.imread('images/cameraman.tif',0)
cv2.imshow("Image read in Python", img)
k = cv2.waitKey(0) & 0xFF
if k == 27: # wait for ESC key to exit
 cv2.destroyAllWindows()

在 Python 中，import命令用于在文件中包含一个库。因此，使用import cv2命令包含cv2库。图像以numpy数组的形式存储，所以numpy也被导入到文件中。imread函数用于以与 C++相同的方式读取图像。所有 OpenCV 函数在 Python 中都必须以cv2.为前缀。imshow函数用于显示图像。在 Python 中，所有 OpenCV 函数都具有与 C++类似的签名和功能。

可以使用以下命令在终端中执行代码：

# For Python2.7
$ python image_read.py
# For Python 3
$ python image_read.py

程序的输出如下所示：

这一节只是为了让您知道 Python 也可以用作编程语言，通过 OpenCV 开发计算机视觉应用，并在 Jetson TX1 上部署它。

摘要

本章描述了在部署 CUDA 和 OpenCV 代码时使用 Jetson TX1。本章详细解释了 TX1 板上 GPU 设备的特性，使其非常适合部署计算复杂的应用。本章测量并比较了 Jetson TX1 在执行如添加两个大型数组这样的 CUDA 应用时的性能，并与书中之前提到的笔记本电脑上的 GPU 进行了比较。本章详细解释了在 Jetson TX1 上处理图像的流程。图像处理应用，如图像相加、图像阈值和图像滤波，在 Jetson TX1 上部署，并对它们的性能进行了测量。

Jetson TX1 的最佳之处在于，可以在嵌入式环境中与多个摄像头进行接口连接，并且可以从该摄像头处理视频以设计复杂的计算机视觉应用。从 Jetson TX1 上连接的板载或 USB 摄像头捕获视频的流程在本文中详细解释。

本章还描述了在 Jetson TX1 上部署高级计算机视觉应用，如人脸检测、眼睛检测和背景减法。Python 语言也可以用于在 Jetson TX1 上部署计算机视觉应用。这一概念在章节的最后部分进行了解释。到目前为止，我们已经看到了如何利用 C/C++ 语言的优势来利用 CUDA 和 GPU 加速。

接下来的几章将演示使用 PyCUDA 模块在 Python 语言中使用 CUDA 和 GPU 加速。

问题

将 Jetson TX1 上的 GPU 设备性能与书中之前提到的 GeForce 940 GPU 进行比较。
判断正误：书中之前提到的所有 CUDA 程序都可以在 Jetson TX1 上执行，无需修改。
在 Jetson TX1 上重新编译 OpenCV 的需要是什么？
判断正误：OpenCV 无法从连接到 USB 端口的摄像头捕获视频。
判断正误：对于计算密集型应用，使用 CSI 摄像头比使用 USB 摄像头更好。
如果你正在使用 OpenCV 开发计算密集型的计算机视觉应用，你更倾向于哪种语言以获得更快的性能？
在 Jetson TX1 上是否有必要安装单独的 OpenCV Python 绑定或 Python 解释器？

第十章：PyCUDA 入门

我们已经看到如何使用 OpenCV 和 CUDA 加速各种应用。我们使用了 C 或 C++ 作为编程语言。如今，Python 在许多领域都非常流行，因此如果我们能够使用 CUDA 加速 Python 应用程序，将会非常有用。Python 提供了一个 PyCUDA 模块，它正是为此而设计的。

它使用 Nvidia CUDA 工具包，这反过来又需要在计算机上安装 Nvidia 显卡。本章将介绍 Python 语言和 PyCUDA 模块，特别是。它将讨论在 Windows 和 Linux 操作系统上安装 PyCUDA 模块的安装过程。尽管本章需要一些对 Python 语言的熟悉，但新来者也能跟随大多数步骤。

本章将涵盖以下主题：

Python 编程语言简介
PyCUDA 模块简介
在 Windows 上安装 PyCUDA
在 Ubuntu 上安装 PyCUDA

技术要求

本章需要良好的 Python 编程语言理解。它还需要任何配备 Nvidia GPU 的计算机或笔记本电脑。本章中使用的 Windows PyCUDA 安装文件可以从以下 GitHub 链接下载：github.com/PacktPublishing/Hands-On-GPU-Accelerated-Computer-Vision-with-OpenCV-and-CUDA。

Python 编程语言简介

Python 正在持续增加其受欢迎程度，因为它可以在许多领域使用，具有广泛的应用。它是一种高级编程语言，可以帮助用几行代码表达复杂的系统。Python 语法易于学习，比其他语言如 C++ 和 Java 更易于阅读，这使得它对新手程序员来说更容易学习。

Python 是一种轻量级的脚本语言，可以轻松用于嵌入式应用。此外，它是一种解释型语言，需要解释器而不是编译器，这与其他编程语言不同。这允许程序员逐行执行代码。它需要一个可以在所有操作系统上轻松安装的 Python 解释器。由于 Python 是开源的，因此一个庞大的社区选择与之合作。他们开发了一系列库，并将其开源，因此它可以无成本地用于应用程序。

Python 可以用于各种领域，如数据科学、机器学习、深度学习、数据分析、图像处理、计算机视觉、数据挖掘和网页开发。它几乎为所有提到的操作系统领域提供了现成的模块，有助于快速开发应用程序。本书前面解释过的 OpenCV 库也具有 Python 接口。因此，它可以轻松地与 Python 代码集成，用于计算机视觉应用程序。Python 还具有用于机器学习和深度学习的库，可以与 OpenCV 一起用于计算机视觉应用程序。

类似于 Python 这样的解释语言的一个缺点是，它比编译语言如 C 或 C++ 慢得多。Python 有一个特性，可以在 Python 脚本中集成 C 或 C++ 代码。这允许你使用 Python 包装器编写计算密集型的 C 或 C++ 代码。

PyCUDA 模块简介

在最后一节中，我们看到了使用 Python 编程语言的优势。还提到 Python 比 C 或 C++ 慢得多。因此，如果它能利用 GPU 的并行处理能力，将会很有益。Python 提供了一个 PyCUDA 包装器，可以通过使用 Nvidia CUDA API 来利用 GPU 的并行计算能力。Python 还有一个 PyOpenCL 模块，可以用于任何 GPU 上的并行计算。

然后，你可能会有一个疑问，为什么你必须使用 PyCUDA，它是专门针对 Nvidia GPU 的。使用 PyCUDA 相比其他类似模块有许多优势；以下是一些原因：

它为 Python 开发者提供了一个与 CUDA API 交互的简单接口，并且有良好的文档，这使得学习变得容易。
可以使用 PyCUDA 模块在 Python 代码中利用 Nvidia 提供的 CUDA API 的全部功能。
PyCUDA 的基础层是用 C++ 编写的，这使得它更快。
它具有更高的抽象级别，与基于 Nvidia C 的运行时 API 相比，使用起来更简单。
它具有非常高效的内存管理机制，与对象的生存期相关联的对象清理功能。这个特性帮助它编写正确的代码，没有内存泄漏或崩溃。
CUDA 代码中的错误也可以通过 Python 异常来处理，这有助于代码中的错误处理机制。

本节描述了使用 PyCUDA 加速 Python 应用程序的优势。在下一节中，我们将看到如何在 Windows 和 Ubuntu 操作系统上安装 PyCUDA 的步骤。

在 Windows 上安装 PyCUDA

本节将描述在 Windows 操作系统上安装 PyCUDA 的步骤。演示使用的是 Windows 10，但该过程适用于任何较新的 Windows 版本。以下是步骤描述：

如果您尚未安装 CUDA 工具包，如第一章所述，请从developer.nvidia.com/cuda-downloads下载最新的 CUDA 工具包。它将询问您的操作系统、CPU 架构以及是否通过互联网安装或先下载整个安装程序。如下面的截图所示，我们选择了带有本地安装程序的 Windows 10。您可以根据您的设置选择相应的值：

双击下载的安装程序并按照屏幕提示安装 CUDA 工具包。
安装具有可视化 C++版本的最新 Visual Studio 版本。我们使用的是免费的 Visual Studio 2017 社区版，其可视化 C++路径应添加到路径环境变量中。可以通过右键单击我的电脑（此电脑）|属性|高级系统设置|环境变量|系统变量来访问环境变量。在路径环境变量中添加可视化 C++安装和 CUDA 工具包安装的 bin 文件夹路径，如下面的截图所示：

将 Anaconda 分布用作 Python 解释器，因此可以从以下网站下载：www.anaconda.com/download/。我们使用的是 Python 3.6 版本的 Anaconda 5.2，如下面的截图所示：

双击下载的安装程序并按照屏幕提示安装 Anaconda。确保勾选将安装路径添加到路径环境变量的复选框。
根据您的系统设置，从以下链接下载最新的 PyCUDA 二进制文件：www.lfd.uci.edu/~gohlke/Pythonlibs/#pycuda。我们使用 CUDA 9.2148 和 Python 3.6，因此选择了相应的 PyCUDA 版本，如下所示：

打开命令提示符，转到 PyCUDA 二进制文件下载的文件夹，并执行以下截图所示的命令：

命令将在 Python 分布中完成 PyCUDA 的安装。

检查 PyCUDA 安装步骤

以下步骤用于检查 PyCUDA 是否已正确安装：

打开 Spyder，这是一个随 Anaconda 安装一起提供的 Python IDE。您可以在开始菜单中输入Spyder来打开它。
在 Spyder IDE 中，如以下截图所示，在 IPython 控制台中输入import pycuda。如果没有错误报告，则表示 PyCUDA 已正确安装。

在 Ubuntu 上安装 PyCUDA

本节将描述在 Linux 操作系统上安装 PyCUDA 的步骤。以 Ubuntu 为演示，但此过程适用于任何最新的 Linux 发行版。步骤如下：

如果您尚未安装 CUDA 工具包，如第一章所述，请从 developer.nvidia.com/cuda-downloads 下载最新的 CUDA 工具包。它将询问您的操作系统、CPU 架构以及是否通过互联网安装或首先下载整个安装程序。如下面的截图所示，我们选择了带有运行文件（本地）安装程序的 Ubuntu。您可以根据您的设置选择值：

在命令提示符中运行 sudo sh cuda_9.2.148_396.37_linux.run 命令以安装 CUDA 工具包。
将使用 Anaconda 发行版作为 Python 解释器，因此可以从网站：www.anaconda.com/download/ 下载和安装。我们使用的是带有 Python 3.6 版本的 Anaconda 5.2，如下面的截图所示：

安装 Anaconda 后，在终端中执行以下命令，如下面的截图所示以安装 PyCUDA：

命令将在 Python 发行版中完成 PyCUDA 的安装。

检查 PyCUDA 安装的步骤

以下步骤用于检查 PyCUDA 是否已正确安装：

打开 Spyder，这是 Anaconda 安装附带的一个 Python IDE。您可以在终端中输入 Spyder 来打开它。
在 Spyder IDE 中，在 IPython 控制台中输入 import pycuda，如下面的截图所示。如果没有错误报告，则表示 PyCUDA 已正确安装。

摘要

总结来说，本章介绍了 Python 编程语言及其在各种领域用于广泛应用的用法。与 C 或 C++ 语言相比，它轻量但较慢。因此，如果它能利用 GPU 的并行计算能力，那么它将非常有用。PyCUDA 是一个 Python 包装器，允许 Python 代码利用 Nvidia CUDA API。PyCUDA 相比于 Python 中可用的其他并行处理模块的优势被详细解释。PyCUDA 使用 Nvidia CUDA 运行时 API 和 Python 解释器。Anaconda 是一个著名的 Python 发行版，它附带了许多有用的 Python 库和 IDE，以及 CUDA 工具包。本章讨论了在 Windows 和 Ubuntu 操作系统上安装 PyCUDA 的详细步骤。

在接下来的两章中，我们将详细介绍如何使用 PyCUDA 加速 Python 应用程序。

问题

Python 相比于 C 或 C++ 等编程语言有哪些优势？
编译型语言和解释型语言之间的区别是什么？
判断对错：Python 比 C 或 C++ 更快。
PyOpenCL 相比 PyCUDA 的优势是什么？
判断对错：Python 允许在 Python 脚本中使用 C 或 C++ 代码。

第十一章：使用 PyCUDA 进行工作

在上一章中，我们看到了为 Windows 和 Linux 操作系统安装 PyCUDA 的步骤。在本章中，我们将首先开发第一个在控制台上显示字符串的 PyCUDA 程序。了解和访问 PyCUDA 运行的 GPU 的设备属性非常重要；这一方法将在本章中详细讨论。我们还将查看 PyCUDA 中内核的线程和块执行。任何 CUDA 编程的重要编程概念，如分配和释放设备上的内存、从主机到设备以及相反的数据传输，以及内核调用，将使用向量加法程序示例进行详细讨论。还将讨论使用 CUDA 事件测量 PyCUDA 程序性能的方法，并将其与 CPU 程序进行比较。这些编程概念将被用于开发一些复杂的 PyCUDA 程序，例如数组元素的平方和矩阵乘法。本章的最后部分描述了在 PyCUDA 中定义内核函数的一些高级方法。

本章将涵盖以下主题：

在 PyCUDA 中编写第一个“Hello, PyCUDA!”程序
从 PyCUDA 程序访问设备属性
PyCUDA 中的线程和块执行
使用向量加法程序的基本 PyCUDA 编程概念
使用 CUDA 事件测量 PyCUDA 程序的性能
PyCUDA 中的一些复杂程序
PyCUDA 中的高级内核函数

技术要求

本章需要良好的 Python 编程语言理解。它还需要任何带有 Nvidia GPU 的计算机或笔记本电脑。本章中使用的所有代码都可以从以下 GitHub 链接下载：github.com/PacktPublishing/Hands-On-GPU-Accelerated-Computer-Vision-with-OpenCV-and-CUDA。

查看以下视频以查看代码的实际运行情况：

bit.ly/2QPWojV

在 PyCUDA 中编写第一个程序

本节描述了使用 PyCUDA 编写简单“Hello, PyCUDA!”程序的步骤。它将演示编写任何 PyCUDA 程序的工作流程。由于 Python 是一种解释型语言，代码也可以从 Python 终端逐行运行，或者保存为.py扩展名并作为文件执行。

使用 PyCUDA 显示从内核中简单字符串的程序如下所示：

import pycuda.driver as drv
import pycuda.autoinit
from pycuda.compiler import SourceModule

mod = SourceModule("""
  #include <stdio.h>

  __global__ void myfirst_kernel()
  {
    printf("Hello,PyCUDA!!!");
  }
""")

function = mod.get_function("myfirst_kernel")
function(block=(1,1,1))

在开发 PyCUDA 代码时的第一步是包含代码所需的所有库。使用import指令来包含一个库、模块、类或函数。这与在 C 或 C++中包含指令类似，并且可以通过以下三种不同的方式完成，如下面的步骤所示。以下也展示了使用三个导入模块的示例：

导入pycuda.driver为drv

这表示导入了 pymodule 的驱动子模块，并给它一个简写符号drv，所以当需要使用pycuda.driver模块中的函数时，可以使用drv.functionname。此模块包含内存管理函数、设备属性、数据方向函数等。
导入pycuda.autoinit

这个命令表示从pycuda的autoint模块导入了。没有给出任何缩写符号。autoint模块用于设备初始化、上下文创建和内存清理。此模块不是必需的，上述所有功能也可以手动完成。
从pycuda.compiler导入SourceModule

这个命令表示只从pycuda.compiler模块导入了SourceModule类。当你只想使用一个大模块中的一个类时，这是很重要的。SourceModule类用于在 PyCUDA 中定义类似 C 的内核函数。

C 或 C++内核代码作为构造函数传递给Sourcemodule类，并创建 mod 对象。内核代码非常简单，因为它只是在控制台上打印一个Hello, PyCUDA!字符串。由于内核代码中使用了printf函数，因此包含stdio.h头文件非常重要。myfirst_kernel函数在内核代码中使用__global__指令定义，表示该函数将在 GPU 上执行。该函数不接受任何参数。它只是在控制台上打印一个字符串。此内核函数将由nvcc编译器编译。

这个函数可以通过使用mod对象的get_function方法创建一个指向该函数的指针在 Python 代码内部使用。内核函数的名称作为引号内的参数给出。指针变量可以取任何名字。这个指针变量用于在代码的最后一行调用内核。内核函数的参数也可以在这里指定，但由于myfirst_kernel函数没有参数，所以没有指定参数。内核可以提供的线程数和每个网格要启动的块数也可以通过使用可选的块和网格参数作为参数提供。块参数被赋予值（1,1,1），这是一个 1 x 3 的 Python 元组，表示块大小为 1 x 1 x 1。因此，将启动一个线程，该线程将在控制台上打印字符串。

程序的输出如下所示：

总结来说，本节逐步展示了开发一个简单的 PyCUDA 程序的步骤。

内核调用

使用 ANSI C 关键字以及 CUDA 扩展关键字编写的设备代码被称为内核。它通过一种名为 内核调用 的方法从 Python 代码中启动。基本上，内核调用的意义是我们从主机代码中启动设备代码。内核代码与正常 C 函数非常相似；只是这段代码是由多个线程并行执行的。它在 Python 中的语法非常简单，如下所示：

kernel (parameters for kernel,block=(tx,ty,tz) , grid=(bx,by,bz))

这从我们要启动的内核函数的指针开始。你应该确保这个内核指针是通过 get_function 方法创建的。然后，它可以包括用逗号分隔的内核函数参数。块参数表示要启动的线程数，而网格参数表示网格中的块数。块和网格参数使用一个 1 x 3 的 Python 元组指定，表示三维空间中的块和线程。内核启动启动的线程总数将是这些数字的乘积。

从 PyCUDA 程序访问 GPU 设备属性

PyCUDA 提供了一个简单的 API 来查找信息，例如，哪些 CUDA 兼容的 GPU 设备（如果有）存在，以及每个设备支持哪些功能。在编写 PyCUDA 程序之前，了解正在使用的 GPU 设备的属性非常重要，这样就可以使用设备的最佳资源。

使用 PyCUDA 显示系统上所有 CUDA 兼容设备属性的程序如下所示：

import pycuda.driver as drv
import pycuda.autoinit
drv.init()
print("%d device(s) found." % drv.Device.count())
for i in range(drv.Device.count()):
  dev = drv.Device(i)
  print("Device #%d: %s" % (i, dev.name()))
  print(" Compute Capability: %d.%d" % dev.compute_capability())
  print(" Total Memory: %s GB" % (dev.total_memory()//(1024*1024*1024)))

  attributes = [(str(prop), value) 
    for prop, value in list(dev.get_attributes().items())]
    attributes.sort()
    n=0

    for prop, value in attributes:
      print(" %s: %s " % (prop, value),end=" ")
      n = n+1
      if(n%2 == 0):
        print(" ")

首先，重要的是要获取系统上存在的 CUDA 兼容设备数量，因为一个系统可能包含多个启用 GPU 的设备。这个数量可以通过 PyCUDA 中驱动类 drv.Device.count() 函数确定。系统上所有设备都会被迭代以确定每个设备的属性。使用 drv.Device 函数为每个设备创建一个指针对象。这个指针用于确定特定设备的所有属性。

name 函数将给出特定设备的名称，而 total_memory 将给出设备上可用的 GPU 全局内存的大小。其他属性存储为 Python 字典，可以通过 get_attributes().items() 函数检索。这通过 Python 中的列表推导式转换为元组列表。这个列表的所有行都包含一个 2 x 1 的元组，其中包含属性的名称和其值。

使用 for 循环迭代此列表以在控制台上显示所有属性。此程序在配备 GeForce 940 GPU 和 CUDA 9 的笔记本电脑上执行。程序输出如下：

这些属性在本书的早期章节中已经详细讨论过，所以我们不再重复讨论；然而，为了总结，本节展示了从 PyCUDA 程序中访问 GPU 设备属性的方法。

PyCUDA 中的线程和块执行

我们在 内核调用 部分看到，我们可以并行启动多个块和多个线程。那么，这些块和线程以什么顺序开始和结束它们的执行？如果我们想在其他线程中使用一个线程的输出，了解这一点很重要。为了理解这一点，我们修改了前面章节中看到的 hello,PyCUDA! 程序中的内核，通过在内核调用中包含一个打印语句来打印块号。修改后的代码如下所示：


import pycuda.driver as drv
import pycuda.autoinit
from pycuda.compiler import SourceModule

mod = SourceModule("""
  #include <stdio.h>
  __global__ void myfirst_kernel()
  {
    printf("I am in block no: %d \\n", blockIdx.x);
  }
""")

function = mod.get_function("myfirst_kernel")
function(grid=(4,1),block=(1,1,1))

从代码中可以看出，我们以并行方式启动了 10 个块，每个块有一个线程。在内核代码中，我们正在打印内核执行的块 ID。我们可以将其视为 10 个相同的 myfirstkernel 并行开始执行。这些副本中的每一个都将有一个唯一的块 ID，可以通过 blockIdx.x 指令访问，以及唯一的线程 ID，可以通过 threadIdx.x 访问。这些 ID 将告诉我们哪个块和线程正在执行内核。当你多次运行程序时，你会发现每次块执行的顺序都不同。一个示例输出如下所示：

它可以产生 n 的阶乘数量的不同输出，其中 n 表示并行启动的块的数量。因此，每次你在 PyCUDA 中编写程序时，都应该小心，因为块会以随机顺序执行。

PyCUDA 中的基本编程概念

在本节中，我们将开始使用 PyCUDA 开发一些有用的功能。本节还将通过一个简单的加法示例展示 PyCUDA 的一些有用的函数和指令。

PyCUDA 中的加法

Python 提供了一个用于数值运算的非常快速的库，称为 numpy (Numeric Python)。它是用 C 或 C++ 开发的，并且对于 Python 中的数组操作非常有用。它在 PyCUDA 程序中经常被用作 PyCUDA 内核函数的参数，这些参数作为 numpy 数组传递。本节解释了如何使用 PyCUDA 添加两个数字。添加两个数字的基本内核代码如下所示：

import pycuda.autoinit
import pycuda.driver as drv
import numpy
from pycuda.compiler import SourceModule
mod = SourceModule("""

  __global__ void add_num(float *d_result, float *d_a, float *d_b)
  {
     const int i = threadIdx.x; 
     d_result[i] = d_a[i] + d_b[i];
  }
""")

如前所述，导入 SourceModule 类和驱动类。numpy 库也被导入，因为它将用于将参数传递给内核代码。add_num 内核函数被定义为 SourceModule 类的构造函数。该函数接受两个设备指针作为输入，一个设备指针指向加法的结果作为输出。需要注意的是，尽管我们在添加两个数字，但内核函数被定义为可以同时处理两个数组加法。两个单个数字不过是每个只有一个元素的数组。如果没有错误，此代码将被编译并加载到设备上。从 Python 调用此内核代码的代码如下所示：


add_num = mod.get_function("add_num")

h_a = numpy.random.randn(1).astype(numpy.float32)
h_b = numpy.random.randn(1).astype(numpy.float32)

h_result = numpy.zeros_like(h_a)
d_a = drv.mem_alloc(h_a.nbytes)
d_b = drv.mem_alloc(h_b.nbytes)
d_result = drv.mem_alloc(h_result.nbytes)
drv.memcpy_htod(d_a,h_a)
drv.memcpy_htod(d_b,h_b)

add_num(
  d_result, d_a, d_b,
  block=(1,1,1), grid=(1,1))
drv.memcpy_dtoh(h_result,d_result)
print("Addition on GPU:")
print(h_a[0],"+", h_b[0] , "=" , h_result[0])

使用 get_function 创建内核函数的指针引用。使用 numpy.random.randn(1) 函数创建两个随机数，该函数用于创建正态分布中的随机数。这些数字使用 astype(numpy.float32) 方法转换为单精度浮点数。用于在主机上存储结果的 numpy 数组被初始化为零。

可以使用 PyCUDA 中驱动类的 mem_alloc 函数在设备上分配内存。内存的大小作为函数的参数传递。使用 h_a.nbytes 函数找到输入的大小。PyCUDA 在驱动类中提供了一个 memcpy 函数，用于从主机内存到设备内存以及相反方向的复制数据。

drv.memcpy_htod 函数用于将数据从主机内存复制到设备内存。设备内存的指针作为第一个参数传递，主机内存指针作为第二个参数传递。通过传递设备指针以及指定要启动的块和线程数量的数字作为参数，调用 add_num 内核。在前面给出的代码中，使用一个线程启动了一个块。内核计算的结果通过使用 drv.memcpy_dtoh 函数复制回主机。结果在控制台上显示，如下所示：

总结来说，本节展示了 PyCUDA 程序的结构。它从内核定义代码开始。然后在 Python 中定义输入。在设备上分配内存并将输入传输到设备内存。接着是内核调用，它将计算结果。然后将结果传输到主机进行进一步处理。PyCUDA 提供了更简单的 API 来执行此操作，这将在下一节中解释。

使用驱动类简化加法程序

PyCUDA 提供了一个更简单的内核调用 API，它不需要内存分配和内存复制。这是通过 API 隐式完成的。这可以通过使用 PyCUDA 中驱动类中的 In 和 Out 函数来实现。修改后的数组加法代码如下所示：

import pycuda.autoinit
import pycuda.driver as drv
import numpy
N = 10
from pycuda.compiler import SourceModule
mod = SourceModule("""

  __global__ void add_num(float *d_result, float *d_a, float *d_b)
 {
    const int i = threadIdx.x; 
    d_result[i] = d_a[i] + d_b[i];
 }
""")
add_num = mod.get_function("add_num")
h_a = numpy.random.randn(N).astype(numpy.float32)
h_b = numpy.random.randn(N).astype(numpy.float32)
h_result = numpy.zeros_like(h_a)
add_num(
  drv.Out(h_result), drv.In(h_a), drv.In(h_b),
  block=(N,1,1), grid=(1,1))
print("Addition on GPU:")
for i in range(0,N):
  print(h_a[i],"+", h_b[i] , "=" , h_result[i])

在前面的代码中，数组中的十个元素被添加，而不是单个元素。内核函数与之前看到的代码完全相同。在主机上创建了两个包含十个随机数的数组。现在，不是创建它们的内存并将它们传输到设备，而是直接调用内核。通过指定数据方向使用 drv.Out 或 drv.In 来修改内核调用。这简化了 PyCUDA 代码并减少了代码的大小。

内核函数使用一个块和每个块 N 个线程。这 N 个线程并行地添加数组中的 N 个元素，从而加速了加法操作。使用 drv.out 指令，内核的结果会自动下载到主机内存，因此这个结果可以直接通过 for 循环打印到控制台。使用 PyCUDA 进行十个元素加法的结果如下所示：

总结来说，本节通过一个简单的数组加法程序介绍了 PyCUDA 的重要概念和函数。使用 PyCUDA 的性能提升可以通过下一节中解释的 CUDA 事件来量化。

使用 CUDA 事件测量 PyCUDA 程序的性能

到目前为止，我们还没有明确确定 PyCUDA 程序的性能。在本节中，我们将看到如何使用 CUDA 事件测量程序的性能。这在 PyCUDA 中是一个非常重要的概念，因为它将允许你从许多选项中选择特定应用的性能最佳算法。

CUDA 事件

我们可以使用 Python 时间测量选项来测量 CUDA 程序的性能，但它不会给出准确的结果。它将包括许多其他因素中的线程延迟在操作系统中的时间开销和调度。使用 CPU 测量的时间也将取决于高精度 CPU 计时器的可用性。很多时候，当 GPU 内核运行时，主机正在执行异步计算，因此 Python 的 CPU 计时器可能不会给出内核执行的正确时间。因此，为了测量 GPU 内核计算的时间，PyCUDA 提供了一个事件 API。

CUDA 事件是在 PyCUDA 程序中指定点记录的 GPU 时间戳。在这个 API 中，GPU 记录时间戳，消除了使用 CPU 计时器测量性能时存在的问题。使用 CUDA 事件测量时间有两个步骤：创建事件和记录事件。我们可以记录两个事件，一个在代码的开始处，一个在结束处。然后我们将尝试计算两个事件之间的时间差，这将给出代码的整体性能。

在 PyCUDA 代码中，可以包含以下行来使用 CUDA 事件 API 测量性能：

import pycuda.driver as drv
start = drv.Event()
end=drv.Event()
#Start Time
start.record()
#The kernel code for which time is to be measured
#End Time
end.record()
end.synchronize()
#Measure time difference
secs = start.time_till(end)*1e-3

使用record方法来测量当前时间戳。在内核代码前后测量时间戳以测量内核执行时间。可以使用time_till方法测量时间戳之间的差异，如前述代码所示。它将以毫秒为单位给出时间，然后将其转换为秒。在下一节中，我们将尝试使用 CUDA 事件来测量代码的性能。

使用大数组加法来测量 PyCUDA 的性能

本节将演示如何使用 CUDA 事件来测量 PyCUDA 程序的性能。同时，还描述了 PyCUDA 代码与简单 Python 代码性能的比较。为了准确比较性能，选取了包含一百万个元素的数组。下面展示了用于大数组加法的内核代码：

import pycuda.autoinit
import pycuda.driver as drv
import numpy
import time
import math

from pycuda.compiler import SourceModule
N = 1000000
mod = SourceModule("""

__global__ void add_num(float *d_result, float *d_a, float *d_b,int N)
{
 int tid = threadIdx.x + blockIdx.x * blockDim.x; 
  while (tid < N)
  {
    d_result[tid] = d_a[tid] + d_b[tid];
    tid = tid + blockDim.x * gridDim.x;
  }
}
""")

由于元素数量较多，因此会启动多个块和线程。所以，既使用线程 ID 又使用块 ID 来计算线程索引。如果启动的总线程数不等于元素数，则同一线程会添加多个元素。这是通过内核函数内部的while循环实现的。同时，它也会确保线程索引不会超出数组元素。除了输入数组和输出数组外，数组的大小也被作为内核函数的参数，因为在SourceModule中内核代码无法访问 Python 的全局变量。下面展示了用于添加大数组的 Python 代码：

start = drv.Event()end=drv.Event()
add_num = mod.get_function("add_num")

h_a = numpy.random.randn(N).astype(numpy.float32)
h_b = numpy.random.randn(N).astype(numpy.float32)

h_result = numpy.zeros_like(h_a)
h_result1 = numpy.zeros_like(h_a)
n_blocks = math.ceil((N/1024))
start.record()
add_num(
  drv.Out(h_result), drv.In(h_a), drv.In(h_b),numpy.uint32(N),
  block=(1024,1,1), grid=(n_blocks,1))
end.record()
end.synchronize()
secs = start.time_till(end)*1e-3
print("Addition of %d element of GPU"%N)
print("%fs" % (secs))

创建了两个事件start和stop来测量 GPU 代码的执行时间。使用驱动类中的Event()函数来定义事件对象。然后，使用get_function创建内核函数的指针引用。使用numpy库的randn函数初始化两个包含一百万个元素的数组，并使用随机数。由于它们是浮点数，因此将它们转换为单精度数以加快设备上的计算速度。

每个块支持 1,024 个线程，正如我们在设备属性部分所看到的。因此，根据这一点，通过将N除以 1,024 来计算总块数。它可能是一个浮点值，因此使用numpy库的ceil函数将其转换为下一个最高整数值。然后，使用计算出的块数和每个块的 1,024 个线程来启动内核。使用numpy.uint32数据类型传递数组的大小。

使用记录函数在调用内核函数前后记录时间，并计算时间差以测量内核函数的执行时间。计算出的时间将打印在控制台上。为了将此性能与 CPU 时间进行比较，程序中添加了以下代码：

start = time.time()
for i in range(0,N):
    h_result1[i] = h_a[i] +h_b[i]
end = time.time()
print("Addition of %d element of CPU"%N)
print(end-start,"s")

使用 Python 的时间库来测量 CPU 时间。使用for循环遍历数组中的每个元素。（注意：由于两个数组都是 numpy 数组，你也可以使用h_result1 = h_a + h_b。）使用time.time()函数在for循环前后测量时间，并将这两个时间之间的差值打印到控制台。程序的输出如下所示：

如输出所示，GPU 添加一百万个元素需要 9.4 毫秒，而 CPU 需要 415.15 毫秒，因此使用 GPU 可以实现大约 50 倍的性能提升。

总结来说，本节展示了使用事件来测量 GPU 代码的计时。将 GPU 性能与 CPU 性能进行比较，以量化使用 GPU 时的性能提升。

PyCUDA 中的复杂程序

到现在为止，PyCUDA 的语法和术语应该已经熟悉了。我们将利用这些知识来开发高级程序，并学习一些 PyCUDA 的高级概念。在本节中，我们将使用 PyCUDA 开发一个程序，使用三种不同的方法对数组的元素进行平方。我们还将学习在 PyCUDA 中进行矩阵乘法的代码。

PyCUDA 中矩阵的逐元素平方

在本节中，使用三种不同的方法执行矩阵中数字的逐元素平方操作。在这个过程中，详细解释了使用多维线程和块的概念、驱动类中的inout指令以及gpuarray类。

使用多维线程的简单内核调用

本节实现了使用 PyCUDA 对矩阵的每个元素进行平方的简单内核函数。以下是一个 5 x 5 矩阵中每个元素平方的内核函数示例：

import pycuda.driver as drv
import pycuda.autoinit 
from pycuda.compiler import SourceModule
import numpy
mod = SourceModule("""
  __global__ void square(float *d_a)
  {
    int idx = threadIdx.x + threadIdx.y*5;
    d_a[idx] = d_a[idx]*d_a[idx];
  }
""")

核函数square只接受一个设备指针作为输入，该指针指向矩阵，并将每个元素替换为其平方。随着多维线程的启动，x和y方向上的线程索引被用来索引矩阵中的值。你可以假设一个 5 x 5 矩阵被展平为一个 1 x 25 的向量，以理解索引机制。请注意，在这个代码中，矩阵的大小是硬编码为5的，但它也可以像上一节中的数组大小一样由用户定义。使用此内核函数的 Python 代码如下所示：

start = drv.Event()
end=drv.Event()
h_a = numpy.random.randint(1,5,(5, 5))
h_a = h_a.astype(numpy.float32)
h_b=h_a.copy()

start.record()

d_a = drv.mem_alloc(h_a.size * h_a.dtype.itemsize)
drv.memcpy_htod(d_a, h_a)

square = mod.get_function("square")
square(d_a, block=(5, 5, 1), grid=(1, 1), shared=0)

h_result = numpy.empty_like(h_a)
drv.memcpy_dtoh(h_result, d_a)
end.record()
end.synchronize()
secs = start.time_till(end)*1e-3
print("Time of Squaring on GPU without inout")
print("%fs" % (secs))
print("original array:")
print(h_a)
print("Square with kernel:")
print(h_result)

创建了两个事件来测量内核函数的计时。在主机上，一个 5x5 的矩阵通过 numpy.random 模块的 randint 函数初始化为随机数。这需要三个参数。前两个参数定义了用于生成随机数的数字范围。第一个参数是最小值，第二个参数是用于生成数字的最大值。第三个参数是大小，指定为元组 (5,5)。这个生成的矩阵再次转换为单精度数以加快处理速度。矩阵的内存是在设备上分配的，生成的随机数矩阵被复制到其中。

创建了指向内核函数的指针引用，并通过传递设备内存指针作为参数调用内核。内核调用使用多维线程，x 和 y 方向上的值为 5。因此，总共启动了 25 个线程，每个线程计算矩阵中单个元素的平方。内核计算的结果被复制回主机并在控制台上显示。内核所需的时间以及输入和输出矩阵都显示在控制台上。

计算一个 5x5 矩阵中每个元素的平方需要 149 毫秒。使用驱动类的 inout 指令可以简化相同的计算。这将在下一节中解释。

使用 `inout` 与内核调用结合使用

如上一节程序的核心函数所示，相同的数组既用作输入也用作输出。PyCUDA 的驱动模块为这类情况提供了一个 inout 指令。它消除了为该数组单独分配内存、上传到设备以及将结果下载回主机的需求。所有操作都在内核调用期间同时进行。这使得代码更简单，更容易阅读。使用驱动类 inout 指令的 Python 代码如下所示：

start.record()
start.synchronize()

square(drv.InOut(h_a), block=(5, 5, 1))

end.record()
end.synchronize()

print("Square with InOut:")
print(h_a)
secs = start.time_till(end)*1e-3
print("Time of Squaring on GPU with inout")
print("%fs" % (secs))

使用 inout 指令初始化 CUDA 事件以测量代码的性能。内核调用与上一节相同，因此在此不再重复。可以看出，在调用平方内核时，通过 drv.inout 指令传递了一个变量作为参数。因此，所有与设备相关的操作都在这一步中完成。内核调用使用多维线程，与上一节的情况相同。计算结果和耗时被打印到控制台，如下所示：

所需时间相对于原始内核来说比较少。因此，通过使用驱动类中的 inout 指令，PyCUDA 代码可以变得高效且易于阅读。PyCUDA 还提供了一个用于数组相关操作的 gpuarray 类。它也可以用于平方操作，这将在下一节中解释。

使用 `gpuarray` 类

Python 提供了一个 numpy 库，用于在 Python 中进行数值计算。PyCUDA 提供了一个与 numpy 类似的 gpuarray 类，该类在 GPU 设备上存储其数据和执行其计算。数组的形状和数据类型与 numpy 中完全相同。gpuarray 类提供了许多用于计算的算术方法。它消除了使用 SourceModule 在 C 或 C++ 中指定内核代码的需要。因此，PyCUDA 代码将只包含 Python 代码。使用 gpuarray 类对矩阵的每个元素进行平方的代码如下所示：

import pycuda.gpuarray as gpuarray
import numpy
import pycuda.driver as drv

start = drv.Event()
end=drv.Event()
start.record()
start.synchronize()

h_b = numpy.random.randint(1,5,(5, 5))
d_b = gpuarray.to_gpu(h_b.astype(numpy.float32))
h_result = (d_b**2).get()
end.record()
end.synchronize()

print("original array:")
print(h_b)
print("doubled with gpuarray:")
print(h_result)
secs = start.time_till(end)*1e-3
print("Time of Squaring on GPU with gpuarray")
print("%fs" % (secs))

gpuarray 类需要在代码中使用，它位于 pycuda.gpuarray 模块中。矩阵使用从 1 到 5 的随机整数进行初始化，以便进行计算。这个矩阵通过使用 gpuarray 类的 to_gpu() 方法上传到设备内存。要上传的矩阵作为参数传递给此方法。矩阵被转换为单精度数字。所有对这个上传矩阵的操作都将在该设备上执行。平方操作以与我们在 Python 代码中执行的方式类似的方式进行，但由于变量是使用 gpuarray 存储在设备上的，因此此操作也将在该设备上执行。结果通过使用 get 方法下载回主机。以下是在控制台上显示的结果，包括使用 gpuarray 进行逐元素平方所需的时间：

计算平方需要大约 58 毫秒。它完全消除了在 C 语言中定义内核函数的需要，其功能与 numpy 库相似，因此 Python 程序员可以轻松地与之一起工作。

总结来说，在本节中，我们使用 PyCUDA 以三种不同的方式开发了一个逐元素平方程序。我们还看到了 PyCUDA 中的多维线程、inout 指令和 gpuarray 类的概念。

使用 GPU 数组进行点积

两个向量之间的点积是各种应用中重要的数学运算。上一节中使用的 gpuarray 类可以用来计算两个向量之间的点积。gpuarray 方法计算点积的性能与 numpy 操作进行了比较。用于使用 numpy 计算点积的代码如下所示：

import pycuda.gpuarray as gpuarray
import pycuda.driver as drv
import numpy
import time
import pycuda.autoinit
n=100
h_a=numpy.float32(numpy.random.randint(1,5,(1,n)))
h_b=numpy.float32(numpy.random.randint(1,5,(1,n)))

start=time.time()
h_result=numpy.sum(h_a*h_b)

#print(numpy.dot(a,b))
end=time.time()-start
print("Answer of Dot Product using numpy")
print(h_result)
print("Time taken for Dot Product using numpy")
print(end,"s")

初始化了两个各有 100 个元素的向量，并使用随机整数来计算点积。Python 的时间模块用于计算计算点积所需的时间。使用*运算符来计算两个向量的逐元素乘积，然后将这些结果相加以计算总的点积。请注意，这里使用的numpy.dot方法用于矩阵乘法，不能用于点积。计算出的点积和时间将在控制台上显示。以下是如何使用gpuarray在 GPU 上执行相同操作的代码：

d_a = gpuarray.to_gpu(h_a)
d_b = gpuarray.to_gpu(h_b)

start1 = drv.Event()
end1=drv.Event()
start1.record()

d_result = gpuarray.dot(d_a,d_b)
end1.record()
end1.synchronize()
secs = start1.time_till(end1)*1e-3
print("Answer of Dot Product on GPU")
print(d_result.get())
print("Time taken for Dot Product on GPU")
print("%fs" % (secs))
if(h_result==d_result.get()):
  print("The computed dor product is correct")

使用to_gpu方法将两个向量上传到 GPU 上以计算点积。gpuarray类提供了一个点积方法，可以直接用于计算点积。它需要一个 GPU 数组作为参数。使用get()方法将计算结果下载回主机。计算结果和通过 CUDA 事件测量的时间将在控制台上显示。程序的结果如下所示：

从输出中可以看出，使用numpy和gpuarray计算点积得到相同的结果。numpy库计算点积需要 37 毫秒，而 GPU 只需 0.1 毫秒即可完成相同的操作。这进一步说明了使用 GPU 和 PyCUDA 进行复杂数学运算的优势。

矩阵乘法

经常使用的一个重要数学运算是矩阵乘法。本节将演示如何使用 PyCUDA 在 GPU 上执行矩阵乘法。当矩阵的大小非常大时，这是一个非常复杂的数学运算。应记住，对于矩阵乘法，第一个矩阵的列数应等于第二个矩阵的行数。矩阵乘法不是累积操作。为了避免复杂性，在这个例子中，我们使用相同大小的方阵。如果你熟悉矩阵乘法的数学，你可能还记得，第一个矩阵的每一行将与第二个矩阵的所有列相乘。这将对第一个矩阵的所有行重复进行。以下是一个 3x3 矩阵乘法的示例：

结果矩阵中的每个元素将通过将第一个矩阵的对应行与第二个矩阵的对应列相乘来计算。这个概念被用来开发以下所示的内核函数：


import numpy as np
from pycuda import driver
from pycuda.compiler import SourceModule
import pycuda.autoinit
MATRIX_SIZE = 3 

matrix_mul_kernel = """
__global__ void Matrix_Mul_Kernel(float *d_a, float *d_b, float *d_c)
{
  int tx = threadIdx.x;
  int ty = threadIdx.y;
  float value = 0;

  for (int i = 0; i < %(MATRIX_SIZE)s; ++i) {
    float d_a_element = d_a[ty * %(MATRIX_SIZE)s + i];
    float d_b_element = d_b[i * %(MATRIX_SIZE)s + tx];
    value += d_a_element * d_b_element;
 }

   d_c[ty * %(MATRIX_SIZE)s + tx] = value;
 } """

matrix_mul = matrix_mul_kernel % {'MATRIX_SIZE': MATRIX_SIZE}

mod = SourceModule(matrix_mul)

内核函数接受两个输入数组和一个输出数组作为参数。矩阵的大小作为常量传递给内核函数。这样就消除了需要将向量的大小作为内核函数参数之一的需求，正如本章前面所解释的那样。两种方法都是正确的，取决于程序员认为哪种更方便。每个线程计算结果矩阵的一个元素。第一矩阵的行和第二矩阵的列的所有元素在for循环内相乘并求和。答案被复制到结果矩阵中的相应位置。内核函数内部计算索引的细节可以在本书的早期章节中找到。以下是如何使用此内核函数的 Python 代码：


h_a = np.random.randint(1,5,(MATRIX_SIZE, MATRIX_SIZE)).astype(np.float32)
h_b = np.random.randint(1,5,(MATRIX_SIZE, MATRIX_SIZE)).astype(np.float32)

d_a = gpuarray.to_gpu(h_a) 
d_b = gpuarray.to_gpu(h_b)
d_c_gpu = gpuarray.empty((MATRIX_SIZE, MATRIX_SIZE), np.float32)

matrixmul = mod.get_function("Matrix_Mul_Kernel")

matrixmul(d_a, d_b,d_c_gpu, 
  block = (MATRIX_SIZE, MATRIX_SIZE, 1),
)
print("*" * 100)
print("Matrix A:")
print(d_a.get())

print("*" * 100)
print("Matrix B:")
print(d_b.get())

print("*" * 100)
print("Matrix C:")
print(d_c_gpu.get())

  # compute on the CPU to verify GPU computation
h_c_cpu = np.dot(h_a, h_b)
if h_c_cpu == d_c_gpu.get() :
    print("The computed matrix multiplication is correct")

两个大小为 3 x 3 的矩阵被初始化为从1到5的随机整数。这些矩阵使用gpuarray类的to_gpu方法上传到设备内存。创建一个空的 GPU 数组以在设备上存储结果。这三个变量作为参数传递给内核函数。内核函数调用时，矩阵大小作为x和y方向的维度。结果使用get()方法下载回主机。两个输入矩阵和 GPU 计算的结果在控制台上打印。使用numpy库的 dot 方法在 CPU 上计算矩阵乘法。结果与内核计算的结果进行比较，以验证内核计算的结果。程序的结果如下所示：

总结来说，我们已经开发了一个简单的内核函数，使用 PyCUDA 执行矩阵乘法。这个内核函数可以通过使用共享内存进一步优化，正如本书前面所解释的那样。

PyCUDA 中的高级内核函数

到目前为止，我们已经看到了使用SourceModule类在 C 或 C++中定义内核函数的使用。我们还使用了gpuarray类来进行设备计算，而不需要显式定义内核函数。本节描述了 PyCUDA 中可用的高级内核定义功能。这些功能用于开发各种并行通信模式的内核函数，如映射、归约和扫描操作。

PyCUDA 中的元素级内核

这个特性允许程序员定义一个内核函数，该函数作用于数组的每个元素。它允许程序员将一个或多个操作数组成的复杂表达式执行为一个单一的计算步骤。以下是这样定义大型数组元素级加法内核函数的方式：

import pycuda.gpuarray as gpuarray
import pycuda.driver as drv
from pycuda.elementwise import ElementwiseKernel
from pycuda.curandom import rand as curand
add = ElementwiseKernel(
  "float *d_a, float *d_b, float *d_c",
  "d_c[i] = d_a[i] + d_b[i]",
  "add")

使用PyCuda.elementwise.ElementwiseKernel函数来定义元素级核函数。它需要三个参数。第一个参数是核函数的参数列表。第二个参数定义了对每个元素要执行的操作，第三个参数指定了核函数的名称。以下是如何使用此核函数的 Python 代码示例：


n = 1000000
d_a = curand(n)
d_b = curand(n)
d_c = gpuarray.empty_like(d_a)
start = drv.Event()
end=drv.Event()
start.record()
add(d_a, d_b, d_c)
end.record()
end.synchronize()
secs = start.time_till(end)*1e-3
print("Addition of %d element of GPU"%shape)
print("%fs" % (secs))
# check the result
if d_c == (d_a + d_b):
  print("The sum computed on GPU is correct")

使用pycuda.curandom类中的curand函数用随机数初始化两个数组。这又是一个有用的功能，因为它消除了在主机上初始化然后上传到设备内存的需要。创建一个空的 GPU 数组来存储结果。通过将这些三个变量作为参数传递来调用add核函数。使用 CUDA 事件计算一百万个元素加法所需的时间，并在控制台上显示。

程序的输出如下所示：

元素级核函数只需要 0.6 毫秒即可完成数组中一百万个元素的加法操作。这种性能优于本章前面看到的程序。因此，当要对向量执行元素级操作时，元素级核定义是一个非常重要的概念需要记住。

归约核

归约操作可以通过使用某些表达式将元素集合归约到单个值来定义。它在各种并行计算应用中非常有用。以计算向量点积的例子来展示 PyCUDA 中的归约概念。以下是如何使用 PyCUDA 中归约核功能计算点积的程序示例：

import pycuda.gpuarray as gpuarray
import pycuda.driver as drv
import numpy
from pycuda.reduction import ReductionKernel
import pycuda.autoinit
n=5
start = drv.Event()
end=drv.Event()
start.record()
d_a = gpuarray.arange(n,dtype= numpy.uint32)
d_b = gpuarray.arange(n,dtype= numpy.uint32)
kernel = ReductionKernel(numpy.uint32,neutral="0",reduce_expr="a+b",map_expr="d_a[i]*d_b[i]",arguments="int *d_a,int *d_b")
d_result = kernel(d_a,d_b).get()
end.record()
end.synchronize()
secs = start.time_till(end)*1e-3
print("Vector A")
print(d_a)
print("Vector B")
print(d_b)
print("The computed dot product using reduction:")
print(d_result)
print("Dot Product on GPU")
print("%fs" % (secs))

PyCUDA 提供了pycuda.reduction.ReductionKernel类来定义归约核。它需要许多参数。第一个参数是输出数据类型。第二个参数是中值，通常定义为0。第三个参数是用于归约元素集合的表达式。在前面代码中定义了加法操作。第四个参数定义为归约前操作数之间映射操作的表达式。在代码中定义了元素级乘法。最后一个参数定义了核函数的参数。

计算点积的归约核函数需要两个向量之间的元素级乘法，然后对所有元素进行加法。使用arange函数定义了两个向量。它在 Python 中的range函数类似，但arange会将数组保存在设备上。通过将这些两个向量作为参数传递来调用核函数，并将结果检索到主机。使用 CUDA 事件计算所需的计算时间，并在控制台上与点积的结果一起显示，如下所示：

减少内核计算点积大约需要 2.5 秒，与上一节中看到的显式内核相比，这是一个相对较长时间。然而，在需要减少操作的并行计算应用中，它非常有用。

扫描内核

扫描操作再次是一个非常重要的并行计算范式。扫描操作符将指定的函数应用于输入的第一个元素。该函数的结果作为输入提供，并带有原始输入的第二个元素。所有中间结果形成输出序列。这个概念可以用于各种应用。以累积加法为例，演示了 PyCUDA 中的扫描内核概念。累积加法不过是将加法应用于向量的每个元素，顺序进行。示例如下：

Input Vector
[7 5 9 2 9]
Scan Operation for cumulative sum
[7,7+5,7+5+9,7+5+9+2,7+2+9+2+7]

如所示，前一次加法的结果被添加到当前元素中，以计算当前位置的输出。这被称为包含扫描操作。如果输入的当前元素不参与，则称为排除扫描。使用包含扫描执行累积求和的程序如下所示：


import pycuda.gpuarray as gpuarray
import pycuda.driver as drv
import numpy
from pycuda.scan import InclusiveScanKernel
import pycuda.autoinit
n=10
start = drv.Event()
end=drv.Event()
start.record()
kernel = InclusiveScanKernel(numpy.uint32,"a+b")
h_a = numpy.random.randint(1,10,n).astype(numpy.int32)
d_a = gpuarray.to_gpu(h_a)
kernel(d_a)
end.record()
end.synchronize()
secs = start.time_till(end)*1e-3
assert(d_a.get() == numpy.cumsum(h_a,axis=0)).all()
print("The input data:")
print(h_a)
print("The computed cumulative sum using Scan:")
print(d_a.get())
print("Cumulative Sum on GPU")
print("%fs" % (secs))

PyCUDA 提供了pycuda.scan.InclusiveScanKernel类来定义一个包含扫描内核。它需要输出数据类型和用于扫描的操作作为参数。对于累积求和，指定了加法操作。随机整数的数组被应用于这个内核函数的输入。内核输出将与输入具有相同的大小。输入和输出向量以及计算累积和所需的时间将在控制台上显示，如下所示：

在一个数组的 10 个元素上运行扫描操作大约需要 2 毫秒。总结来说，在本节中，我们看到了定义映射、减少和扫描操作内核的各种特殊方法。

摘要

本章展示了 PyCUDA 编程的概念。它从使用 PyCUDA 开发一个简单的 Hello, PyCUDA 程序开始。详细讨论了在 C 或 C++ 中定义内核以及在 Python 代码中调用它的概念，以及从 PyCUDA 程序中访问 GPU 设备属性的 API。通过一个简单的程序解释了 PyCUDA 程序中多线程和多块执行机制。使用一个数组加法的简单示例描述了 PyCUDA 程序的基本结构。通过驱动类指令描述了 PyCUDA 代码的简化。详细解释了使用 CUDA 事件来衡量 PyCUDA 程序性能的方法。使用逐元素平方示例解释了驱动类中的 inout 指令和 gpuarray 类的功能。使用 gpuarray 类开发了使用 PyCUDA 计算点积的代码。详细解释了 PyCUDA 中用于矩阵乘法等复杂数学运算的代码。本章的最后部分描述了用于映射、归约和扫描操作的多种内核定义方法。

下一章将在此基础上构建知识，并描述 PyCUDA 中可用的某些高级内核以及使用 PyCUDA 开发计算机视觉应用程序。

问题

在 PyCUDA 中使用 SourceModule 类定义内核函数时，使用哪种编程语言？将使用哪种编译器来编译这个内核函数？
为本章中使用的 myfirst_kernel 函数编写一个内核调用函数，块的数量等于 1024 x 1024，每个块中的线程数等于 512 x 512。
判断对错：PyCUDA 程序内部的块执行是按顺序进行的。
在 PyCUDA 程序中使用 In、Out 和 inout 驱动类原语的优势是什么？
编写一个 PyCUDA 程序，使用 gpuarray 类将向量中每个元素的值增加 2，向量的大小是任意的。
使用 CUDA 事件来测量内核执行时间的优势是什么？
判断对错：gpuarray 类是 Python 中 numpy 库的 GPU 设备版本。

第十二章：使用 PyCUDA 的基本计算机视觉应用

在上一章中，我们学习了与 PyCUDA 相关的重要编程概念。我们还学习了如何使用这些编程概念在 PyCUDA 中开发一些程序。本章将在此基础上构建知识，我们将使用 PyCUDA 来开发基本的图像处理和计算机视觉应用。原子操作和共享内存的并行编程概念也将被详细解释。图像直方图传达了与图像对比度相关的信息，它还可以用作计算机视觉任务的图像特征。本章将详细解释使用 PyCUDA 计算直方图的程序。其他基本的计算机视觉应用，如使用 PyCUDA 进行颜色空间转换、图像加法和图像反转，也将被描述。

本章将涵盖以下主题：

使用原子操作和共享内存进行直方图计算
使用 PyCUDA 的基本计算机视觉应用
从网络摄像头进行图像和视频的颜色空间转换
图像加法
图像反转

技术要求

本章要求对 Python 编程语言有良好的理解。它还需要任何带有 Nvidia GPU 的计算机或笔记本电脑。本章中使用的所有代码都可以从以下 GitHub 链接下载：github.com/PacktPublishing/Hands-On-GPU-Accelerated-Computer-Vision-with-OpenCV-and-CUDA。查看以下视频以查看代码的实际运行情况：bit.ly/2prC1wI

PyCUDA 中的直方图计算

图像直方图传达了与图像对比度相关的信息，它还可以用作计算机视觉任务的图像特征。直方图表示特定像素值出现的频率。在计算大小为 256 x 256 的 8 位图像的直方图时，65,535 个像素值将在 0-255 的强度值数组上工作。如果为每个像素启动一个线程，那么 65,535 个线程将在 256 个强度值内存位置上工作。

考虑这样一种情况，即大量线程试图修改内存的一小部分。在计算图像直方图时，必须对所有内存位置执行读取-修改-写入操作。这个操作是 d_out[i] ++，其中首先从内存中读取 d_out[i]，然后增加，最后写回内存。然而，当多个线程在相同的内存位置执行此操作时，可能会给出错误的结果。

假设一个内存位置有一个初始值0，线程a和b尝试递增这个内存位置，那么最终答案应该是2。然而，在执行时，可能会发生a和b线程同时读取这个值的情况，那么这两个线程都将得到值0。它们将其递增到1，并且两个线程都将这个1存储在内存中。所以，计算出的答案是1，这是不正确的。

为了理解这可能会多么危险，考虑 ATM 现金取款的例子。假设你在账户中有$50,000 的余额。你有两张同一账户的 ATM 卡。你和你的朋友同时去两个不同的 ATM 取款$40,000。你们同时刷卡；所以，当 ATM 检查余额时，两个都会得到$50,000。如果你俩都取款$40,000，那么两台机器都会查看初始余额，即$50,000。取款金额小于余额，因此两台机器都会给出$40,000。尽管你的余额是$50,000，但你得到了$80,000，这是危险的。为了避免这些情况，在并行编程中使用原子操作，这将在下一节中解释。

使用原子操作

CUDA 提供了一个名为atomicAdd的 API，用于避免并行访问内存位置时的问题。这是一个阻塞操作，这意味着当多个线程尝试访问相同的内存位置时，一次只能有一个线程可以访问该内存位置。其他线程必须等待此线程完成并将答案写入内存。以下是一个使用atomicAdd操作计算直方图的内核函数示例：

import pycuda.autoinit
import pycuda.driver as drv
import numpy
import matplotlib.pyplot as plt
from pycuda.compiler import SourceModule
mod = SourceModule("""          
__global__ void atomic_hist(int *d_b, int *d_a, int SIZE)
{
 int tid = threadIdx.x + blockDim.x * blockIdx.x;
 int item = d_a[tid];
 if (tid < SIZE)
 {
  atomicAdd(&(d_b[item]), 1);
 }
}
""")

内核函数有三个参数。第一个参数是计算后存储直方图输出的输出数组。对于 8 位图像，这个数组的大小将是 256。第二个参数是图像强度的展平数组。第三个参数是展平数组的大小。在线程索引处按像素强度索引的直方图数组的内存位置将为每个线程递增。线程的数量等于展平图像数组的大小。

atomicAdd函数用于递增内存位置。它接受两个参数。第一个是我们想要递增的内存位置，第二个是这个位置需要递增的值。atomicadd函数将增加直方图计算在执行时间上的成本。使用原子操作计算直方图的 Python 代码如下：

atomic_hist = mod.get_function("atomic_hist")
import cv2
h_img = cv2.imread("cameraman.tif",0)

h_a=h_img.flatten()
h_a=h_a.astype(numpy.int)
h_result = numpy.zeros(256).astype(numpy.int)
SIZE = h_img.size
NUM_BIN=256
n_threads= int(numpy.ceil((SIZE+NUM_BIN-1) / NUM_BIN))
start = drv.Event()
end=drv.Event()
start.record()
atomic_hist(
    drv.Out(h_result), drv.In(h_a), numpy.uint32(SIZE),
    block=(n_threads,1,1), grid=(NUM_BIN,1))

end.record()
end.synchronize()
secs = start.time_till(end)*1e-3
print("Time for Calculating Histogram on GPU with shared memory")
print("%fs" % (secs)) 
plt.stem(h_result)
plt.xlim([0,256])
plt.title("Histogram on GPU")

使用get_function()方法创建指向内核函数的指针引用。图像读取使用 OpenCV 库。如果 Python 中没有安装，您可以从命令提示符执行以下命令：

$pip install opencv-python

然后，可以使用import cv2命令从任何 Python 程序中导入OpenCV库。图像读取函数与本书前面解释过的类似。图像被读取为灰度图像。在 Python 中，图像被存储为一个numpy数组。这个数组被展平为一个向量，以便它可以被一维线程和块操作。也可以在不展平的情况下使用二维线程处理图像。numpy库提供了一个flatten()方法来执行此操作。

块和线程的总数是根据图像的大小和直方图的桶数来计算的。在调用核函数时，将展平的图像数组、空直方图数组和展平数组的大小作为参数传递，同时传递要启动的块和线程数。核函数返回计算出的直方图，可以显示或绘制。

Python 提供了一个包含丰富绘图函数的matplotlib库。该库中的stem函数用于绘制离散的直方图函数。xlim函数用于设置 X 轴的界限。title函数用于给图表添加标题。程序的输出如下所示：

如果直方图没有所有强度的均匀分布，则可能导致对比度差的图像。可以通过直方图均衡化来增强对比度，这会将这种分布转换为均匀分布。直方图还传达了关于图像亮度的信息。如果直方图集中在图表的左侧，则图像会太暗；如果集中在右侧，则图像会太亮。再次强调，可以使用直方图均衡化来纠正这个问题。

计算直方图的核函数也可以使用并行编程中共享内存的概念来开发。这将在以下部分中说明。

使用共享内存

共享内存可以在 GPU 设备上片上使用，因此它比全局内存快得多。共享内存的延迟大约是全球未缓存内存延迟的 100 倍。来自同一块的线程都可以访问共享内存。这在许多需要线程之间共享结果的场景中非常有用。然而，如果不进行同步，这也可能导致混乱或错误的结果。如果一个线程在另一个线程写入之前从内存中读取数据，可能会导致错误的结果。因此，这种内存访问应该得到适当的控制或管理。这通过__syncthreads()指令来完成，它确保在程序继续前进之前，所有对内存的写入操作都已完成。这也被称为屏障。屏障的含义是所有线程都将到达这一行并等待其他线程完成。当所有线程都到达这个屏障后，它们可以继续前进。本节将演示如何从 PyCUDA 程序中使用共享内存。

这种共享内存的概念可以用于计算图像的直方图。内核函数如下所示：

import pycuda.autoinit
import pycuda.driver as drv
import numpy
import matplotlib.pyplot as plt
from pycuda.compiler import SourceModule

mod1 = SourceModule("""
__global__ void atomic_hist(int *d_b, int *d_a, int SIZE)
{
 int tid = threadIdx.x + blockDim.x * blockIdx.x;
 int offset = blockDim.x * gridDim.x;
 __shared__ int cache[256];
 cache[threadIdx.x] = 0;
 __syncthreads();

 while (tid < SIZE)
 {
  atomicAdd(&(cache[d_a[tid]]), 1);
  tid += offset;
 }
 __syncthreads();
 atomicAdd(&(d_b[threadIdx.x]), cache[threadIdx.x]);
}
""")

对于 8 位图像，bins 的数量是 256，所以我们定义的共享内存大小等于块中的线程数，也就是 bins 的数量。我们将为当前块计算一个直方图，因此共享内存被初始化为零，并按前面讨论的方式计算该块的直方图。但这次，结果存储在共享内存中而不是全局内存中。在这种情况下，只有 256 个线程试图访问共享内存中的 256 个内存元素，而不是之前代码中的所有 65,535 个元素。这将有助于减少原子操作中的开销时间。最后一行的最终原子加操作将一个块的直方图添加到整体直方图值中。由于加法是一个累积操作，我们不必担心每个块执行的顺序。以下是如何使用此内核函数计算直方图的 Python 代码示例：

atomic_hist = mod.get_function("atomic_hist")

import cv2
h_img = cv2.imread("cameraman.tif",0)

h_a=h_img.flatten()
h_a=h_a.astype(numpy.int)
h_result = numpy.zeros(256).astype(numpy.int)
SIZE = h_img.size
NUM_BIN=256
n_threads= int(numpy.ceil((SIZE+NUM_BIN-1) / NUM_BIN))
start = drv.Event()
end=drv.Event()
start.record()
atomic_hist(
 drv.Out(h_result), drv.In(h_a), numpy.uint32(SIZE),
 block=(n_threads,1,1), grid=(NUM_BIN,1),shared= 256*4)

end.record()
end.synchronize()
secs = start.time_till(end)*1e-3
print("Time for Calculating Histogram on GPU with shared memory")
print("%fs" % (secs)) 
plt.stem(h_result)
plt.xlim([0,256])
plt.title("Histogram on GPU")

代码几乎与上一节相同。唯一的区别在于内核调用。在调用内核时应该定义共享内存的大小。这可以通过内核调用函数中的共享参数来指定。它被指定为256*4，因为共享内存的大小为 256 个整数元素，每个元素需要 4 字节存储。与上一节显示的相同直方图将被显示。

为了检查计算出的直方图的准确性并比较性能，还使用 OpenCV 内置函数calcHist计算直方图，如下面的代码所示：

start = cv2.getTickCount()
hist = cv2.calcHist([h_img],[0],None,[256],[0,256])
end = cv2.getTickCount()
time = (end - start)/ cv2.getTickFrequency()
print("Time for Calculating Histogram on CPU")
print("%fs" % (secs))

calcHist 函数需要五个参数。第一个参数是图像变量的名称。第二个参数在彩色图像的情况下指定通道。对于灰度图像，它为零。第三个参数指定了如果您想计算图像特定部分的直方图，则需要指定掩码。第四个参数指定了箱子的数量，第五个参数指定了强度值的范围。OpenCV 还在 Python 中提供了 getTickCount 和 getTickFrequency 函数来计算 OpenCV 代码的性能。没有共享内存、使用共享内存以及使用 OpenCV 函数的性能如下所示：

没有共享内存的内核函数耗时 1 毫秒，而使用共享内存时为 0.8 毫秒，这进一步证明了使用共享内存可以提高内核函数性能的观点。总结来说，在本节中，我们看到了在 GPU 上计算直方图的两种不同方法。我们还了解了原子操作和共享内存的概念，以及如何在 PyCUDA 中使用它们。

使用 PyCUDA 的基本计算机视觉操作

本节将演示如何使用 PyCUDA 开发简单的计算机视觉应用。在 Python 中，图像不过是二维或三维的 numpy 数组，因此在 PyCUDA 中处理和操作图像与处理多维数组类似。本节将为您提供一个基本的概念，帮助您开发一个简单的应用，您可以用它来利用 PyCUDA 开发复杂的计算机视觉应用。

PyCUDA 中的颜色空间转换

大多数计算机视觉算法都处理灰度图像，因此需要将相机捕获的彩色图像转换为灰度图像。尽管 OpenCV 提供了内置函数来完成此操作，但您也可以通过开发自己的函数来实现。本节将演示如何开发一个 PyCUDA 函数，用于将彩色图像转换为灰度图像。如果已知将图像从一个颜色空间转换为另一个颜色空间的公式，那么本节中显示的函数可以通过替换公式来为任何颜色空间转换编写。

OpenCV 以 BGR 格式捕获和存储图像，其中蓝色是第一个通道，接着是绿色和红色。从 BGR 格式转换为灰度的公式如下：

gray = 0.299*r+0.587*g+0.114*b Where r,g,b are color intensities of red, green and blue channel at a particular location

该函数在图像和视频中的实现将在以下两个子节中展示。

图像上的 BGR 到灰度转换

在本节中，我们将尝试开发将 BGR 图像转换为灰度图像的内核函数。将彩色图像转换为灰度的内核函数如下所示：

import pycuda.driver as drv
from pycuda.compiler import SourceModule
import numpy as np
import cv2
mod = SourceModule \
  (
    """
#include<stdio.h>
#define INDEX(a, b) a*256+b

__global__ void bgr2gray(float *d_result,float *b_img, float *g_img, float *r_img)
{
 unsigned int idx = threadIdx.x+(blockIdx.x*(blockDim.x*blockDim.y));
 unsigned int a = idx/256;
 unsigned int b = idx%256;
 d_result[INDEX(a, b)] = (0.299*r_img[INDEX(a, b)]+0.587*g_img[INDEX(a, b)]+0.114*b_img[INDEX(a, b)]);

}
 """
)

定义了一个小的 INDEX 函数来计算一个 256 x 256 大小的二维图像的特定索引值。将彩色图像三个通道的展平图像数组作为核函数的输入，其输出是相同大小的灰度图像。INDEX 函数用于将线程索引转换为图像中的特定像素位置。使用所示函数计算该位置的灰度值。以下是将彩色图像转换为灰度图像的 Python 代码示例：

h_img = cv2.imread('lena_color.tif',1)
h_gray=cv2.cvtColor(h_img,cv2.COLOR_BGR2GRAY)
#print a
b_img = h_img[:, :, 0].reshape(65536).astype(np.float32)
g_img = h_img[:, :, 1].reshape(65536).astype(np.float32)
r_img = h_img[:, :, 2].reshape(65536).astype(np.float32)
h_result=r_img
bgr2gray = mod.get_function("bgr2gray")
bgr2gray(drv.Out(h_result), drv.In(b_img), drv.In(g_img),drv.In(r_img),block=(1024, 1, 1), grid=(64, 1, 1))

h_result=np.reshape(h_result,(256,256)).astype(np.uint8)
cv2.imshow("Grayscale Image",h_result)
cv2.waitKey(0)
cv2.destroyAllWindows()

使用 OpenCV 的 imread 函数读取彩色图像。图像的大小应为 256 x 256，因此如果不是这个大小，则应使用 cv2.resize 函数将其转换为该大小。彩色图像以 BGR 格式存储，因此使用 Python 的数组切片将其蓝色、绿色和红色通道分离。这些数组被展平，以便可以将它们传递给核函数。

核函数使用三个颜色通道作为输入和一个用于存储输出灰度图像的数组进行调用。核函数将在每个像素位置计算灰度值，并返回一个灰度图像的展平数组。使用 numpy 库的 reshape 函数将此结果数组转换回原始图像大小。OpenCV 的 imshow 函数需要无符号整数数据类型来显示图像，因此数组也转换为 uint8 数据类型。灰度图像显示在屏幕上，如下面的截图所示：

在摄像头视频中执行 BGR 到灰度的转换

在上一节中开发的将图像转换为灰度的相同核函数可以用来将来自摄像头的视频转换为灰度。以下是其 Python 代码示例：

cap = cv2.VideoCapture(0)
bgr2gray = mod.get_function("bgr2gray")
while(True):
  # Capture frame-by-frame
  ret, h_img = cap.read()
  h_img = cv2.resize(h_img,(256,256),interpolation = cv2.INTER_CUBIC)

  b_img = h_img[:, :, 0].reshape(65536).astype(np.float32)
  g_img = h_img[:, :, 1].reshape(65536).astype(np.float32)
  r_img = h_img[:, :, 2].reshape(65536).astype(np.float32)
  h_result=r_img

  bgr2gray(drv.Out(h_result), drv.In(b_img), drv.In(g_img),drv.In(r_img),block=(1024, 1, 1), grid=(64, 1, 1))

  h_result=np.reshape(h_result,(256,256)).astype(np.uint8)
  cv2.imshow("Grayscale Image",h_result)

  # Display the resulting frame
  cv2.imshow('Original frame',h_img)
  if cv2.waitKey(50) & 0xFF == ord('q'):
    break

# When everything done, release the capture
cap.release()
cv2.destroyAllWindows()

Python 中的 OpenCV 提供了一个 VideoCapture 类来从摄像头捕获视频。它需要一个相机设备索引作为参数。对于摄像头，它被指定为零。然后，启动一个连续的 while 循环来从摄像头捕获帧。使用捕获对象的 read 方法读取这些帧。这些帧使用 cv2 库的 resize 函数调整大小到 256 x 256。这些帧是彩色图像，因此从它们中分离出三个通道并展平，以便可以将它们传递给核函数。以与上一节相同的方式调用核函数，并将结果重塑以在屏幕上显示。以下是一帧摄像头流的代码输出：

网络摄像头流将一直持续到在键盘上按下 q 键。总结一下，我们已经在 PyCUDA 中开发了一个内核函数，用于将 BGR 格式的彩色图像转换为灰度图像，这个函数既可以处理图像也可以处理视频。这些内核函数可以通过替换相同的方程来修改，以用于其他颜色空间的转换。

PyCUDA 中的图像加法

当两个图像大小相同时，可以执行两个图像的加法。它执行两个图像的逐像素加法。假设在两个图像中，(0,0) 像素的强度值分别为 100 和 150，那么结果图像中的强度值将是 250，这是两个强度值的和，如下所示：

result = img1 + img2

OpenCV 加法是一个饱和操作，这意味着如果加法的结果超过 255，则它将被饱和在 255。因此，相同的函数性作为 PyCUDA 内核函数实现。执行图像加法的代码如下所示：

import pycuda.driver as drv
from pycuda.compiler import SourceModule
import numpy as np
import cv2
mod = SourceModule \
 (
"""
 __global__ void add_num(float *d_result, float *d_a, float *d_b,int N)
{
 int tid = threadIdx.x + blockIdx.x * blockDim.x; 
 while (tid < N)
  {
 d_result[tid] = d_a[tid] + d_b[tid];
 if(d_result[tid]>255)
 {
 d_result[tid]=255;
 }
 tid = tid + blockDim.x * gridDim.x;
}
}
"""
)
img1 = cv2.imread('cameraman.tif',0)
img2 = cv2.imread('circles.png',0)
h_img1 = img1.reshape(65536).astype(np.float32)
h_img2 = img2.reshape(65536).astype(np.float32)
N = h_img1.size
h_result=h_img1
add_img = mod.get_function("add_num")
add_img(drv.Out(h_result), drv.In(h_img1), drv.In(h_img2),np.uint32(N),block=(1024, 1, 1), grid=(64, 1, 1))
h_result=np.reshape(h_result,(256,256)).astype(np.uint8)
cv2.imshow("Image after addition",h_result)
cv2.waitKey(0)
cv2.destroyAllWindows()

内核函数与上一章中看到的数组加法内核函数类似。内核函数中添加了饱和条件，表示如果像素强度在加法后超过 255，则它将被饱和在 255。读取两个大小相同的图像，展平，并转换为单精度浮点数据类型。这些展平的图像及其大小作为参数传递给内核函数。内核函数计算的结果被重塑为原始图像大小，并使用 imshow 函数转换为无符号整型以显示。结果如下面的截图所示，包括原始图像：

同样的内核函数经过轻微修改后，可以用于其他算术和逻辑运算。

使用 gpuarray 在 PyCUDA 中进行图像反转

除了算术运算外，NOT 运算也广泛用于图像反转，其中黑色转换为白色，白色转换为黑色。它可以表示为以下方程：

result_image = 255 - input_image

在前面的方程中，255 表示 8 位图像的最大强度值。PyCUDA 提供的 gpuarray 类用于开发图像反转程序，如下所示：

import pycuda.driver as drv
import numpy as np
import cv2
import pycuda.gpuarray as gpuarray
import pycuda.autoinit

img = cv2.imread('circles.png',0)
h_img = img.reshape(65536).astype(np.float32)
d_img = gpuarray.to_gpu(h_img)
d_result = 255- d_img
h_result = d_result.get()
h_result=np.reshape(h_result,(256,256)).astype(np.uint8)
cv2.imshow("Image after addition",h_result)
cv2.waitKey(0)
cv2.destroyAllWindows()

图像被读取为灰度图像，展平并转换为单精度浮点数据类型以进行进一步处理。它使用 gpuarray 类的 to_gpu 方法上传到 GPU。使用前面的方程在 GPU 上执行反转，然后使用 get() 方法将结果下载回主机。结果通过重塑为原始图像大小显示在屏幕上，如下面的截图所示：

总结，本节展示了 PyCUDA 在开发基本计算机视觉操作中的应用，如颜色空间转换、图像加法和图像反转。这个概念可以用于使用 PyCUDA 开发复杂的计算机视觉应用。

摘要

本章描述了在开发简单的计算机视觉应用中使用 PyCUDA。它描述了使用 PyCUDA 计算数组直方图的过程。直方图是图像的一个非常重要的统计全局特征，可以用来获取关于图像的重要信息。本章以直方图计算为例，详细解释了原子操作和共享内存的概念。Python 中的图像存储为numpy数组，因此在 PyCUDA 中操作图像类似于修改多维numpy数组。本章描述了 PyCUDA 在多种基本计算机视觉应用中的使用，例如图像加法、图像反转和颜色空间转换。本章中描述的概念可以用于使用 PyCUDA 开发复杂的计算机视觉应用。

本章也标志着本书的结束，本书描述了使用 CUDA 编程和 GPU 硬件加速计算机视觉应用。

问题

判断对错：使用d_out[i]++行而不是atomicadd操作将在直方图计算中产生准确的结果。
使用原子操作共享内存的优势是什么？
当在内核中使用共享内存时，内核调用函数中的修改是什么？
通过计算图像的直方图可以获得哪些信息？
判断对错：本章开发的将 BGR 转换为灰度的内核函数也将适用于 RGB 转换为灰度。
为什么本章中展示的所有示例都将图像展平？这是一个强制性的步骤吗？
为什么在显示之前将图像从numpy库的uint8数据类型转换？

第十三章：评估

第一章

提高性能的三种选项如下：
- 拥有更快的时钟速度
- 单个处理器每个时钟周期完成更多的工作
- 许多可以并行工作的小型处理器。这个选项被 GPU 用来提高性能。
正确
CPU 的设计是为了提高延迟，而 GPU 的设计是为了提高吞吐量。
汽车需要 4 小时才能到达目的地，但它只能容纳 5 人，而可以容纳 40 人的公交车需要 6 小时才能到达目的地。公交车每小时可以运输 6.66 人，而汽车每小时可以运输 1.2 人。因此，汽车具有更好的延迟，而公交车具有更好的吞吐量。
图像不过是一个二维数组。大多数计算机视觉应用都涉及这些二维数组的处理。这涉及到对大量数据进行类似操作，这些操作可以通过 GPU 高效地执行。因此，GPU 和 CUDA 在计算机视觉应用中非常有用。
错误
printf语句在主机上执行

第二章

通过传递参数作为值来减去两个数字的 CUDA 程序如下：

include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>
__global__ void gpuSub(int d_a, int d_b, int *d_c) 
{
 *d_c = d_a - d_b;
}
int main(void) 
{
  int h_c;
  int *d_c;
  cudaMalloc((void**)&d_c, sizeof(int));
 gpuSub << <1, 1 >> > (4, 1, d_c);
 cudaMemcpy(&h_c, d_c, sizeof(int), cudaMemcpyDeviceToHost);
 printf("4-1 = %d\n", h_c);
 cudaFree(d_c);
 return 0;
}

通过传递参数作为引用来乘以两个数字的 CUDA 程序如下：

#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>
 __global__ void gpuMul(int *d_a, int *d_b, int *d_c) 
{
 *d_c = *d_a * *d_b;
}
int main(void) 
{
 int h_a,h_b, h_c;
 int *d_a,*d_b,*d_c;
 h_a = 1;
 h_b = 4;
 cudaMalloc((void**)&d_a, sizeof(int));
 cudaMalloc((void**)&d_b, sizeof(int));
 cudaMalloc((void**)&d_c, sizeof(int));
 cudaMemcpy(d_a, &h_a, sizeof(int), cudaMemcpyHostToDevice);
 cudaMemcpy(d_b, &h_b, sizeof(int), cudaMemcpyHostToDevice);
 gpuMul << <1, 1 >> > (d_a, d_b, d_c);
 cudaMemcpy(&h_c, d_c, sizeof(int), cudaMemcpyDeviceToHost);
 printf("Passing Parameter by Reference Output: %d + %d = %d\n", h_a, h_b, h_c);
 cudaFree(d_a);
 cudaFree(d_b);
 cudaFree(d_c);
 return 0;
 }

为gpuMul内核启动 5000 个线程的三种方法如下：

1\. gpuMul << <25, 200 >> > (d_a, d_b, d_c);
2\. gpuMul << <50, 100 >> > (d_a, d_b, d_c);
3\. gpuMul << <10, 500 >> > (d_a, d_b, d_c);

错误
查找具有 5.0 或更高版本的 GPU 设备的程序如下

int main(void) 
{ 
  int device; 
  cudaDeviceProp device_property; 
  cudaGetDevice(&device); 
  printf("ID of device: %d\n", device); 
  memset(&device_property, 0, sizeof(cudaDeviceProp)); 
  device_property.major = 5; 
  device_property.minor = 0; 
  cudaChooseDevice(&device, &device_property); 
  printf("ID of device which supports double precision is: %d\n", device);                                                                         
  cudaSetDevice(device); 
}

查找数字立方的 CUDA 程序如下：

#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
#define N 50
__global__ void gpuCube(float *d_in, float *d_out) 
{
     //Getting thread index for current kernel
     int tid = threadIdx.x; // handle the data at this index
     float temp = d_in[tid];
     d_out[tid] = temp*temp*temp;
 }
int main(void) 
{
     float h_in[N], h_out[N];
     float *d_in, *d_out;
     cudaMalloc((void**)&d_in, N * sizeof(float));
     cudaMalloc((void**)&d_out, N * sizeof(float));
      for (int i = 0; i < N; i++) 
    {
         h_in[i] = i;
     }
   cudaMemcpy(d_in, h_in, N * sizeof(float), cudaMemcpyHostToDevice);
   gpuSquare << <1, N >> >(d_in, d_out);
  cudaMemcpy(h_out, d_out, N * sizeof(float), cudaMemcpyDeviceToHost);
    printf("Cube of Number on GPU \n");
     for (int i = 0; i < N; i++) 
     {
         printf("The cube of %f is %f\n", h_in[i], h_out[i]);
     }
     cudaFree(d_in);
     cudaFree(d_out);
     return 0;
 }

特定应用的通信模式如下所示：
1. 图像处理 - 模板
2. 移动平均 - 聚合
3. 按升序排序数组 - 散射
4. 在数组中查找数字的立方 - 映射

第三章

选择线程数和块数的最佳方法如下：

gpuAdd << <512, 512 >> >(d_a, d_b, d_c);

每个块可以启动的线程数量有限，对于最新的处理器来说，这个数量是 512 或 1024。同样，每个网格的块数量也有限制。因此，如果有大量线程，那么最好通过少量块和线程来启动内核，如上所述。

以下是为查找 50000 个数字的立方而编写的 CUDA 程序：

#include "stdio.h"
#include<iostream>
#include <cuda.h>
#include <cuda_runtime.h>
#define N 50000
__global__ void gpuCube(float *d_in, float *d_out) 
{
      int tid = threadIdx.x + blockIdx.x * blockDim.x; 
while (tid < N)
{
    float temp = d_in[tid];
    d_out[tid] = temp*temp*temp;
    tid += blockDim.x * gridDim.x;
 }
}
int main(void) 
{
     float h_in[N], h_out[N];
     float *d_in, *d_out;
     cudaMalloc((void**)&d_in, N * sizeof(float));
     cudaMalloc((void**)&d_out, N * sizeof(float));
      for (int i = 0; i < N; i++) 
    {
         h_in[i] = i;
     }
   cudaMemcpy(d_in, h_in, N * sizeof(float), cudaMemcpyHostToDevice);
   gpuSquare << <512, 512 >> >(d_in, d_out);
  cudaMemcpy(h_out, d_out, N * sizeof(float), cudaMemcpyDeviceToHost);
    printf("Cube of Number on GPU \n");
     for (int i = 0; i < N; i++) 
     {
         printf("The cube of %f is %f\n", h_in[i], h_out[i]);
     }
     cudaFree(d_in);
     cudaFree(d_out);
     return 0;
 }

正确，因为它只需要访问局部内存，这是一种更快的内存。
当内核的变量不适合寄存器文件时，它们使用局部内存。这被称为寄存器溢出。因为一些数据不在寄存器中，它将需要更多时间从内存中检索它。这将花费更多时间，因此程序的性能将受到影响。
不，因为所有线程都在并行运行。所以数据可能在写入之前就被读取，因此可能不会给出期望的输出。
正确。在原子操作中，当一个线程正在访问特定的内存位置时，其他所有线程都必须等待。当许多线程访问相同的内存位置时，这将产生时间开销。因此，原子操作会增加 CUDA 程序的执行时间。
Stencil 通信模式非常适合纹理内存。
当在 if 语句中使用 __syncthreads 指令时，对于具有此条件的线程，false 永远不会到达这个点，__syncthreads 将持续等待所有线程到达这个点。因此，程序将永远不会终止。

第四章

CPU 计时器将包括操作系统中的线程延迟和调度的时间开销，以及其他许多因素。使用 CPU 测量的时间也将取决于高精度 CPU 计时器的可用性。主机在 GPU 内核运行时经常执行异步计算，因此 CPU 计时器可能无法给出内核执行的准确时间。
从 C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v9.0\libnvvp 打开 Nvidia Visual profiler。然后，转到 -> 新会话并选择矩阵乘法示例的 .exe 文件。您可以可视化您代码的性能。
除以零、变量类型或大小不正确、不存在变量、下标超出范围等是语义错误的例子。
可以给出线程发散的例子如下：

__global__ void gpuCube(float *d_in, float *d_out) 
{
     int tid = threadIdx.x; 
if(tid%2 == 0)
{
     float temp = d_in[tid];
     d_out[tid] = temp*temp*temp;
 }
else
{
    float temp = d_in[tid];
    d_out[tid] = temp*temp*temp;
}
}

在代码中，奇数和偶数线程执行不同的操作，因此它们完成所需的时间不同。在 if 语句之后，这些线程将再次合并。这将产生时间开销，因为快速线程必须等待慢速线程。这将降低代码的性能。

cudaHostAlloc 函数应谨慎使用，因为这种内存不会被交换到磁盘上；您的系统可能会耗尽内存。这可能会影响系统上运行的其他应用程序的性能。
在 CUDA 流操作中，操作顺序很重要，因为我们希望重叠内存复制操作与内核执行操作。因此，操作队列应设置为这些操作可以相互重叠，否则使用 CUDA 流不会提高程序的性能。
对于 1024 x 1024 的图像，线程数应为 32x32（如果您的系统支持每个块 1024 个线程），块数也应为 32 x 32，这可以通过将图像大小除以每个块线程数来确定。

第五章

图像处理和计算机视觉领域之间存在差异。图像处理关注通过修改像素值来提高图像的视觉质量，而计算机视觉关注从图像中提取重要信息。因此，在图像处理中，输入和输出都是图像，而在计算机视觉中，输入是图像，但输出是从该图像中提取的信息。
OpenCV 库在 C、C++、Java 和 Python 语言中都有接口，并且可以在 Windows、Linux、Mac 和 Android 等所有操作系统上使用，而无需修改单行代码。这个库还可以利用多核处理。它可以利用 OpenGL 和 CUDA 进行并行处理。由于 OpenCV 轻量级，它也可以在树莓派等嵌入式平台上使用。这使得它在实际场景中部署计算机视觉应用成为理想选择。
初始化图像为红色的命令如下：

 Mat img3(1960,1960, CV_64FC3, Scalar(0,0,255) )

从网络摄像头捕获视频并将其存储在磁盘上的程序如下：

#include <opencv2/opencv.hpp>
#include <iostream>

using namespace cv;
using namespace std;

int main(int argc, char* argv[])
{
   VideoCapture cap(0); 
   if (cap.isOpened() == false) 
   {
     cout << "Cannot open Webcam" << endl;
     return -1;
 }
  Size frame_size(640, 640);
  int frames_per_second = 30;

  VideoWriter v_writer("images/video.avi", VideoWriter::fourcc('M', 'J', 'P', 'G'), frames_per_second, frame_size, true); 
  cout<<"Press Q to Quit" <<endl;
  String win_name = "Webcam Video";
  namedWindow(win_name); //create a window
   while (true)
   {
     Mat frame;
     bool flag = cap.read(frame); // read a new frame from video 
     imshow(win_name, frame);
     v_writer.write(frame);
  if (waitKey(1) == 'q')
  {
     v_writer.release(); 
     break;
  }
 }
return 0;
}

OpenCV 使用 BGR 颜色格式来读取和显示图像。
从网络摄像头捕获视频并将其转换为灰度的程序如下：

#include <opencv2/opencv.hpp>
#include <iostream>

using namespace cv;
using namespace std;

int main(int argc, char* argv[])
{
   VideoCapture cap(0); 
 if (cap.isOpened() == false) 
 {
    cout << "Cannot open Webcam" << endl;
    return -1;
 }
 cout<<"Press Q to Quit" <<endl;
 String win_name = "Webcam Video";
 namedWindow(win_name); //create a window
 while (true)
 {
    Mat frame;
    bool flag = cap.read(frame); // read a new frame from video 
    cvtColor(frame, frame,cv::COLOR_BGR2GRAY);
    imshow(win_name, frame);
  if (waitKey(1) == 'q')
  {
      break;
  }
 }
return 0;
}

测量加法和减法操作性能的 OpenCV 程序如下：

#include <iostream>
#include "opencv2/opencv.hpp"

int main (int argc, char* argv[])
{
    //Read Two Images 
    cv::Mat h_img1 = cv::imread("images/cameraman.tif");
    cv::Mat h_img2 = cv::imread("images/circles.png");
    //Create Memory for storing Images on device
    cv::cuda::GpuMat d_result1,d_result2,d_img1, d_img2;
    cv::Mat h_result1,h_result2;
int64 work_begin = getTickCount(); 
    //Upload Images to device     
    d_img1.upload(h_img1);
    d_img2.upload(h_img2);

    cv::cuda::add(d_img1,d_img2, d_result1);
    cv::cuda::subtract(d_img1, d_img2,d_result2);
    //Download Result back to host
    d_result1.download(h_result1);
     d_result2.download(h_result2);
    int64 delta = getTickCount() - work_begin;
//Frequency of timer
    double freq = getTickFrequency();
    double work_fps = freq / delta;
    std::cout<<"Performance of Thresholding on CPU: " <<std::endl;
    std::cout <<"Time: " << (1/work_fps) <<std::endl;   
    cv::waitKey();
    return 0;
}

OpenCV 程序用于执行位运算 AND 和 OR 操作如下：

include <iostream>
#include "opencv2/opencv.hpp"

int main (int argc, char* argv[])
{
    cv::Mat h_img1 = cv::imread("images/cameraman.tif");
    cv::Mat h_img2 = cv::imread("images/circles.png");
    cv::cuda::GpuMat d_result1,d_result2,d_img1, d_img2;
    cv::Mat h_result1,h_result2;
    d_img1.upload(h_img1);
    d_img2.upload(h_img2);

    cv::cuda::bitwise_and(d_img1,d_img2, d_result1);
    cv::cuda::biwise_or(d_img1, d_img2,d_result2);

    d_result1.download(h_result1);
     d_result2.download(h_result2);
cv::imshow("Image1 ", h_img1);
    cv::imshow("Image2 ", h_img2);
    cv::imshow("Result AND operation ", h_result1);
cv::imshow("Result OR operation ", h_result2);
    cv::waitKey();
    return 0;
}

第六章

打印任何颜色图像在(200,200)位置像素强度的 OpenCV 函数如下：

cv::Mat h_img2 = cv::imread("images/autumn.tif",1);
cv::Vec3b intensity1 = h_img1.at<cv::Vec3b>(cv::Point(200, 200));
std::cout<<"Pixel Intensity of color Image at (200,200) is:" << intensity1 << std::endl;

使用双线性插值方法将图像调整大小到(300,200)像素的 OpenCV 函数如下：

cv::cuda::resize(d_img1,d_result1,cv::Size(300, 200), cv::INTER_LINEAR);

使用AREA插值将图像上采样 2 倍的 OpenCV 函数如下：

int width= d_img1.cols;
int height = d_img1.size().height;
cv::cuda::resize(d_img1,d_result2,cv::Size(2*width, 2*height), cv::INTER_AREA);

错误。随着滤波器大小的增加，模糊程度也会增加。
错误。中值滤波器不能去除高斯噪声。它可以去除椒盐噪声。
在应用拉普拉斯算子以去除噪声敏感性之前，必须使用平均或高斯滤波器对图像进行模糊处理。
实现顶帽和黑帽形态学操作的 OpenCV 函数如下：

cv::Mat element = cv::getStructuringElement(cv::MORPH_RECT,cv::Size(5,5)); 
  d_img1.upload(h_img1);
  cv::Ptr<cv::cuda::Filter> filtert,filterb;
  filtert = cv::cuda::createMorphologyFilter(cv::MORPH_TOPHAT,CV_8UC1,element);
  filtert->apply(d_img1, d_resulte);
  filterb = cv::cuda::createMorphologyFilter(cv::MORPH_BLACKHAT,CV_8UC1,element);
  filterb->apply(d_img1, d_resultd);

第七章

从视频中检测黄色物体的 OpenCV 代码如下：请注意，这里没有重复样板代码。

cuda::cvtColor(d_frame, d_frame_hsv, COLOR_BGR2HSV);

//Split HSV 3 channels
cuda::split(d_frame_hsv, d_frame_shsv);

//Threshold HSV channels for Yellow color
cuda::threshold(d_frame_shsv[0], d_thresc[0], 20, 30, THRESH_BINARY);
cuda::threshold(d_frame_shsv[1], d_thresc[1], 100, 255, THRESH_BINARY);
cuda::threshold(d_frame_shsv[2], d_thresc[2], 100, 255, THRESH_BINARY);

//Bitwise AND the channels
cv::cuda::bitwise_and(d_thresc[0], d_thresc[1],d_intermediate);
cv::cuda::bitwise_and(d_intermediate, d_thresc[2], d_result);
d_result.download(h_result);
imshow("Thresholded Image", h_result); 
imshow("Original", frame);

当物体的颜色与背景颜色相同时，基于颜色的目标检测将失败。即使有光照变化，也可能失败。
Canny 边缘检测算法的第一步是高斯模糊，这可以去除图像中存在的噪声。之后，计算梯度。因此，检测到的边缘将比之前看到的其他边缘检测算法受噪声影响更小。
当图像受到高斯或椒盐噪声的影响时，霍夫变换的结果非常差。为了改善结果，必须在预处理步骤中通过高斯和中值滤波器对图像进行滤波。
当计算 FAST 关键点的强度阈值较低时，则更多的关键点将通过段测试并被分类为关键点。随着这个阈值的增加，检测到的关键点数量将逐渐减少。
在 SURF 中，Hessian 阈值的较大值将导致更少但更显著的特征点，而较小值将导致更多但不太显著的特征点。
当 Haar 级联的缩放因子从 1.01 增加到 1.05 时，图像大小在每一尺度上都会以更大的因子减小。因此，每帧需要处理的图像更少，这减少了计算时间；然而，这可能导致无法检测到某些对象。
MoG 相比于 GMG 算法在背景减法方面更快且噪声更少。可以将开闭等形态学操作应用于 GMG 的输出，以减少存在的噪声。

第八章

Jetson TX1 在每秒 Tera 级浮点运算性能方面优于 Raspberry Pi。因此，Jetson TX1 可以用于计算密集型应用，如计算机视觉和深度学习，以实现实时部署。
Jetson TX1 开发板支持多达六个 2 通道或三个 4 通道相机。它附带一个 500 万像素的相机。
必须使用 USB 集线器来连接 Jetson TX1 与超过两个 USB 外设。
True
False. Jetson TX1 包含一个 1.73 GHz 运行的 ARM Cortex A57 四核 CPU。
尽管 Jetson TX1 预装了 Ubuntu，但它不包含计算机视觉应用所需的任何软件包。Jetpack 包含 Tegra (L4T) 板支持包的 Linux，TensorRT，用于计算机视觉应用中的深度学习推理，最新的 CUDA 工具包，cuDNN，这是 CUDA 深度神经网络库，Visionworks，也用于计算机视觉和深度学习应用，以及 OpenCV。因此，通过安装 Jetpack，我们可以快速安装构建计算机视觉应用所需的全部软件包。

第九章

Jetson TX1 上的 GPU 设备全局内存大约为 4 GB，GPU 时钟速度约为 1 GHz。这个时钟速度比本书之前使用的 GeForce 940 GPU 慢。内存时钟速度仅为 13 MHz，而 GeForce 940 为 2.505 GHz，这使得 Jetson TX1 更慢。L2 缓存为 256 KB，而 GeForce 940 为 1 MB。大多数其他属性与 GeForce 940 相似。
True
在最新的 Jetpack 中，OpenCV 没有编译 CUDA 支持，也没有 GStreamer 支持，这是从代码中访问相机所需的。因此，移除 Jetpack 中包含的 OpenCV 安装，并使用 CUGA 和 GStreamer 支持编译新的 OpenCV 版本是个好主意。
False. OpenCV 可以从连接到 Jetson TX1 板的 USB 和 CSI 相机捕获视频。
True. CSI 相机更接近硬件，因此读取帧的速度比 USB 相机快，因此在计算密集型应用中最好使用 CSI 相机。
Python OpenCV 绑定不支持 CUDA 加速，因此对于计算密集型任务，最好使用 C++ OpenCV 绑定。
No. Jetson TX1 预装了 python2 和 python3 解释器，同时 OpenCV 也为 Jetson TX1 编译了；它还安装了 python 二进制文件，因此无需单独安装 python OpenCV 绑定。

第十章

Python 是开源的，拥有庞大的用户社区，他们通过模块为语言做出贡献。这些模块可以轻松地用少量代码在短时间内开发应用程序。Python 语言的语法易于阅读和解释，这使得它对新程序员来说更容易学习。它是一种允许逐行执行代码的解释型语言。这些都是 Python 相对于 C/C++ 的几个优点。
在编译型语言中，整个代码被检查并转换为机器代码，而在解释型语言中，每次只翻译一条语句。解释型语言分析源代码所需的时间较少，但与编译型语言相比，整体执行时间较慢。解释型语言不会像编译型语言那样生成中间代码。
错误。Python 是一种解释型语言，这使得它比 C/C++ 慢。
PyOpenCL 可以利用任何图形处理单元，而 PyCUDA 需要 Nvidia GPU 和 CUDA 工具包。
正确。Python 允许在 Python 脚本中包含 C/C++ 代码，因此计算密集型任务可以写成 C/C++ 代码以实现更快的处理，并为它创建 Python 包装器。PyCUDA 可以利用这一功能来处理内核代码。

第十一章

C/C++ 编程语言用于在 SourceModule 类中编写内核函数，并且这个内核函数由 nvcc（Nvidia C）编译器编译。
内核调用函数如下：

myfirst_kernel(block=(512,512,1),grid=(1024,1014,1))

错误。在 PyCUDA 程序中，块执行的顺序是随机的，PyCUDA 程序员无法确定。
驱动类中的指令消除了为数组单独分配内存、将其上传到设备以及将结果下载回主机的要求。所有操作都在内核调用期间同时执行。这使得代码更简单，更容易阅读。
在数组中每个元素加二的 PyCUDA 代码如下所示：

import pycuda.gpuarray as gpuarray
import numpy
import pycuda.driver as drv

start = drv.Event()
end=drv.Event()
start.record()
start.synchronize()
n=10
h_b = numpy.random.randint(1,5,(1,n))
d_b = gpuarray.to_gpu(h_b.astype(numpy.float32))
h_result = (d_b + 2).get()
end.record()
end.synchronize()

print("original array:")
print(h_b)
print("doubled with gpuarray:")
print(h_result)
secs = start.time_till(end)*1e-3
print("Time of adding 2 on GPU with gpuarray")
print("%fs" % (secs))

使用 Python 时间测量选项来测量 PyCUDA 程序的性能不会给出准确的结果。它将包括许多其他因素中的线程延迟在操作系统中的时间开销和调度。使用时间类测量的时间也将取决于高精度 CPU 定时器的可用性。很多时候，主机在进行异步计算的同时 GPU 内核正在运行，因此 Python 的 CPU 计时器可能无法给出内核执行的正确时间。我们可以通过使用 CUDA 事件来克服这些缺点。
正确

第十二章

错误。这一行代表一个读取-修改-写入操作，当多个线程试图增加相同的内存位置时，如直方图计算的情况，可能会产生错误的结果。
在使用共享内存的情况下，较少的线程试图访问共享内存中的 256 个内存元素，而不是没有共享内存时所有线程的情况。这将有助于减少原子操作中的时间开销。
在使用共享内存的情况下，内核调用函数如下：

atomic_hist(
        drv.Out(h_result), drv.In(h_a), numpy.uint32(SIZE),
        block=(n_threads,1,1), grid=(NUM_BIN,1),shared= 256*4)

在调用内核时，应该定义共享内存的大小。这可以通过在内核调用函数中使用共享参数来指定。

直方图是一种统计特征，它提供了关于图像对比度和亮度的关键信息。如果它具有均匀分布，那么图像将具有良好的对比度。直方图还传达了关于图像亮度的信息。如果直方图集中在图表的左侧，那么图像将太暗，如果集中在右侧，那么图像将太亮。
真的。因为 RGB 和 BGR 颜色格式相同，只是通道的顺序不同。转换的方程式仍然保持不变。
与多维线程和块相比，处理单维线程和块更简单。它简化了内核函数内部的索引机制，因此在每个章节中出现的示例中都进行了这种简化。如果我们正在处理多维线程和块，则这不是强制性的。
imshow函数，用于在屏幕上显示图像，需要一个无符号整数的图像。因此，在屏幕上显示之前，所有由内核函数计算的结果都转换为numpy库的uint8数据类型。

posted @ 2025-09-21 12:12 绝不原创的飞龙阅读(147) 评论(0) 收藏举报

刷新页面返回顶部

龙哥盟

掠夺·扩张·投机·博弈

OpenCV-和-CUDA-GPU-加速的计算机视觉实用指南-全-

OpenCV 和 CUDA GPU 加速的计算机视觉实用指南（全）

前言

本书面向的对象

本书涵盖的内容

要充分利用本书

下载示例代码文件

下载彩色图像

代码实战

使用的约定

联系我们

评论

第一章：介绍 CUDA 并开始使用 CUDA

技术要求

介绍 CUDA

并行处理

介绍 GPU 架构和 CUDA

CUDA 架构

CUDA 应用

CUDA 开发环境

支持 CUDA 的 GPU

NVIDIA 显卡驱动程序

标准 C 编译器

CUDA 开发套件

在所有操作系统上安装 CUDA 工具包

Windows

Linux

Mac

CUDA C 中的基本程序

在 Windows 上创建 CUDA C 程序的步骤

在 Ubuntu 上创建 CUDA C 程序的步骤

摘要

问题

第二章：使用 CUDA C 进行并行编程

技术要求

CUDA 程序结构

CUDA C 中的双变量加法程序

内核调用

配置内核参数

CUDA API 函数

向 CUDA 函数传递参数

按值传递参数

按引用传递参数

在设备上执行线程

从 CUDA 程序中访问 GPU 设备属性

通用设备属性

与内存相关的属性

与线程相关的属性

CUDA 中的向量运算

两个向量加法程序

比较 CPU 和 GPU 代码之间的延迟

CUDA 中向量的逐元素平方

并行通信模式

映射

Gather

散列

模板

转置

摘要

问题

第三章：线程、同步和内存

技术要求

线程

内存架构

全局内存

局部内存和寄存器

缓存内存

线程同步

共享内存

原子操作

常量内存

纹理内存

点积和矩阵乘法示例

点积

矩阵乘法

摘要

问题

第四章：CUDA 的高级概念