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并注册，以便将文件直接通过电子邮件发送给您。

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

1. 在www.packt.com登录或注册。

2. 选择“支持”选项卡。

3. 点击“代码下载与勘误”。

4. 在搜索框中输入书籍名称，并遵循屏幕上的说明。

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

• 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 程序的基本步骤如下：

1. 在主机和设备内存中为数据分配内存。

2. 将数据从主机内存复制到设备内存。

3. 通过指定并行度来启动内核。

4. 在所有线程完成后，将数据从设备内存复制回主机内存。

5. 释放主机和设备上使用的所有内存。

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 的步骤，如下所示：

1. 双击安装程序。它将要求您选择临时安装文件将被提取的文件夹。选择您选择的文件夹。建议保持默认设置。

2. 然后，安装程序将检查系统兼容性。如果您的系统兼容，您可以按照屏幕提示安装 CUDA。您可以选择快速安装（默认）或自定义安装。自定义安装允许您选择要安装的 CUDA 功能。建议选择快速默认安装。

3. 安装程序还将安装 CUDA 示例程序和 CUDA Visual Studio 集成。

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

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

1. 如果您选择了默认安装路径，所有 CUDA 示例都将位于 C:\ProgramData\NVIDIA Corporation\CUDA Samples\vx.x。

2. 要检查安装，您可以运行任何项目。

3. 我们使用位于 C:\ProgramData\NVIDIA Corporation\CUDA Samples\vx.x\1_Utilities\deviceQuery 的设备查询项目。

4. 双击您 Visual Studio 版本的 *.sln 文件。它将在 Visual Studio 中打开此项目。

5. 然后，您可以在 Visual Studio 中点击本地 Windows 调试器。如果构建成功并显示以下输出，则表示安装完成：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/a44f83ca-c438-4b21-8901-493ff5c0a47a.png

Linux

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

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

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

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

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

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

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

sudo apt-get update

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

sudo apt-get install cuda

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

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

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

1. 设置 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 的输出应类似于以下内容：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/08ce73ed-45be-4f7a-bf0c-1da1060bb208.png

Mac

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

1. 启动安装程序，按照屏幕上的提示完成安装。它将安装所有先决条件、CUDA、工具包和 CUDA 示例。

2. 然后，您需要使用以下命令设置环境变量，以指向 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}}

1. 运行脚本：cuda-install-samples-x.x.sh。它将以写权限安装 CUDA 示例。

2. 完成后，您可以去 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 程序的步骤。步骤如下：

1. 打开 Microsoft Visual Studio。

2. 前往文件 | 新建 | 工程。

3. 选择 NVIDIA | CUDA 9.0 | CUDA 9.0 Runtime。

4. 给项目命名你想要的名称，然后点击确定。

5. 它将创建一个包含示例 kernel.cu 文件的工程。现在通过双击它来打开此文件。

6. 从文件中删除现有的代码，并写入之前给出的代码。

7. 从“构建”选项卡构建项目，并按 Ctrl + F5 调试代码。如果一切正常，你将在命令行看到 Hello, CUDA! 如此显示：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/08c42110-ea31-453a-8396-80935ab8d035.png

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

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

1. 通过在终端中打开终端并输入 nsight 来打开 Nsight。

2. 前往文件 | 新建 | CUDA C/C++ 项目。

3. 给项目起一个你喜欢的名字，然后点击确定。

4. 它将创建一个带有示例文件的项目。现在通过双击它来打开此文件。

5. 从文件中删除现有的代码，并写入之前给出的代码。

6. 通过按播放按钮运行代码。如果一切正常，你将在终端看到 Hello, CUDA! 如此显示：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/3ae3b999-ffe4-4963-9239-b44031d6e601.png

摘要

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

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

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

问题

1. 解释三种提高计算硬件性能的方法。哪种方法用于开发 GPU？

2. 正误判断：提高延迟将提高吞吐量。

3. 填空：CPU 设计用于提高 ___，而 GPU 设计用于提高 ___。

4. 以从一个地方到另一个地方旅行为例，距离为 240 公里。你可以选择一辆可以容纳五人的汽车，速度为 60 公里/小时，或者一辆可以容纳 40 人的公共汽车，速度为 40 公里/小时。哪种选项将提供更好的延迟，哪种选项将提供更好的吞吐量？

5. 解释使 GPU 和 CUDA 在计算机视觉应用中特别有用的原因。

6. 正误判断：CUDA 编译器不能编译没有设备代码的代码。

7. 在本章讨论的 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__ 关键字定义的，因此它旨在在设备上执行，而主函数将在主机上执行。该程序在设备上添加两个变量，并在命令行上打印输出，如下所示：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/9869b1ed-27f0-432f-ab68-0b41b03af8e1.png

在本书中，我们将使用一个约定，即主机变量将以前缀 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 复制回主机内存，但这次数据传输方向是从设备到主机。程序输出如下：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/9097c292-8aee-4dbc-964a-31ec2a203559.png

程序结束时使用 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 将告诉我们哪个块和线程正在执行内核。当你多次运行程序时，你会发现，每次块执行的顺序都不同。一个示例输出如下所示：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/50a8685b-96c7-4229-8ac2-2562c57a9a02.png

您应该问的一个问题是，前面的程序将产生多少种不同的输出模式？正确的答案是 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 上执行并组合所有先前代码段输出的结果：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/71c144c7-d4f6-46e1-bd68-3558fd3778f1.png

你可能会问的一个问题是，为什么你应该对了解设备属性感兴趣。答案是，这将帮助你在存在多个 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 函数将答案存储在第三个数组中。然后，我们在控制台上打印这个答案，如下所示：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/783fe03c-7410-4177-9649-d6aa378983ea.png

使用 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 函数释放分配给三个设备指针的内存。程序的输出如下：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/2e6cda1b-587e-4e29-b716-4459ad2ddfe2.png

所有 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 并同时执行时，输出如下：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/085e74b0-e73e-485d-bf4d-1a3dc7f3ffc9.png

从输出中可以看出，当相同的函数在 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 个线程的单个块。程序输出如下：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/f31562a9-874e-4188-9432-a6b379875f16.png

每次你使用这种方式并行启动 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 中相互通信。如果一个线程需要其他线程的数据，那么可以做什么也没有讨论。因此，在下一章中，我们将详细讨论内存架构和线程同步。

问题

1. 编写一个 CUDA 程序来减去两个数字。在内核函数中通过值传递参数。

2. 编写一个 CUDA 程序来乘以两个数字。在内核函数中通过引用传递参数。

3. 假设你想要并行启动 5,000 个线程。以三种不同的方式配置内核参数来完成此操作。每个块最多可以有 512 个线程。

4. 对或错：程序员可以决定在设备上块将按何种顺序执行，以及块将被分配到哪个流多处理器？

5. 编写一个 CUDA 程序以找出你的系统包含一个主次版本为 5.0 或更高版本的 GPU 设备。

6. 编写一个 CUDA 程序来找到一个包含从 0 到 49 的数字的向量的立方。

7. 对于以下应用，哪种通信模式是有用的？

   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 计算的变化是由于并行启动多个线程和块。可以通过将块和线程视为一个二维矩阵来理解这个计算，其中块的数量等于行数，列数等于每个块的线程数。以下是一个例子，有三个块和三个线程/块，如下表所示：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/af773f86-ee8e-4395-af5f-b6139b230bac.png

我们可以通过使用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 值较大的问题。我们不再打印一个非常长的向量的加法，而只打印一条指示计算结果是否正确的打印语句。代码的输出将如下所示：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/e869db0f-e17a-4b05-84b4-9d28f024085d.png

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

内存架构

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

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/ec8e0d7d-05a0-4eca-9c5a-648df0f49a97.png

如图中所示，每个线程都有自己的本地内存和寄存器文件。与处理器不同，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 的值填充这个内存块。然后将其复制回主机内存以打印。结果如下所示：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/a67ad366-a6c8-4cc6-bd8d-3b8e7327152a.png

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

局部内存和寄存器

局部内存和寄存器文件是每个线程独有的。寄存器文件是每个线程可用的最快内存。当内核的变量不适合寄存器文件时，它们会使用局部内存。这被称为寄存器溢出。基本上，局部内存是全局内存的一部分，对每个线程来说是唯一的。与寄存器文件相比，访问局部内存会较慢。尽管局部内存被缓存在 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变量将属于每个线程，并存储在寄存器文件中。当这个变量在内核函数中进行计算时，计算将是最快的。前述代码的输出如下所示：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/01037099-35aa-46bd-bf0b-3ffc7eb92540.png

缓存内存

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

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/d25f4dfd-77ed-4a21-8a9b-3ec1b3a7ca9c.png

本节演示了当多个线程使用同一内存位置的数据时共享内存的使用。下一节将演示使用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 个线程。内核执行后在设备上存储的答案被复制回主机，每个内存位置的价值在控制台上显示。

输出如下：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/220c69ea-0945-4d74-9c6c-58b48f265382.png

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

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/4492405e-802e-4611-8571-0b5e79f7c4e4.png

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

为了解决这个问题，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 次。以下将是输出结果：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/29bea1ee-443f-46cb-9742-0841c6ff3b47.png

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

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

常量内存

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指令的第二次调用中省略了它们。

代码的输出如下所示：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/e604f1e3-6ad4-4b42-a22e-55ad71959484.png

有一个需要注意的事项是常量内存是只读内存。这个例子只是用来解释从 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 数组使用的内存。程序在控制台上显示的输出如下：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/047e2343-482f-480a-aa49-834805450ca6.png

本节结束了我们对 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)。我们打印出数学计算出的点积答案，以及是否计算正确。最后释放主机和设备内存。程序的输出如下：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/ec2ca87b-9356-443f-83ee-d4d81cdcc920.png

矩阵乘法

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

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/64671ce5-f20d-4b13-a9d2-b0ac3767f5c6.png

同样的数据被多次重用，因此这是使用共享内存的理想情况。在本节中，我们将制作两个分别使用和不使用共享内存的单独的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。

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

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/5f13d0c1-02da-4efa-8f76-3da8da6f4735.png

线性数组的索引可以通过其行 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)

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

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/b86bd951-bfd0-4b27-96b2-a20663044300.png

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

摘要

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

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

问题

1. 假设你想要并行启动 100,000 个线程。在网格中块的数量和每个块中的线程数量最佳选择是什么，为什么？

2. 编写一个 CUDA 程序，找出数组中每个元素的立方值，当数组中的元素数量为 100,000 时。

3. 判断以下陈述是对还是错，并给出理由：局部变量之间的赋值运算符将比全局变量之间的赋值运算符更快。

4. 注册溢出是什么？它如何损害你的 CUDA 程序的性能？

5. 判断以下代码行是否会给出所需输出：d_out[i] = d_out[i-1]。

6. 判断以下陈述是对还是错，并给出理由：原子操作会增加 CUDA 程序的执行时间。

7. 在你的 CUDA 程序中使用纹理内存的理想通信模式有哪些？

8. 在 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。它有三个参数。第一个是我们想要存储差异的变量，第二个是开始事件，第三个是结束事件。计算完这个时间后，我们将在下一行将其打印到控制台。我们将此性能测量代码添加到上一章中看到的向量加法代码中，使用了多个线程和块。添加这些行后的输出如下：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/07822ae7-a24a-4f50-b732-a4468731e6d8.png

在 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 视觉分析工具，如下所示：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/fb1a410f-a1dd-4536-bec1-67a77db0610b.png

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

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/ac336aad-6c71-4343-a7ef-fb8ae154a36f.png

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

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/dfaaa8a7-df80-4009-9164-ac73eebdee5a.png

可以看到，当我们从主机复制两个数组到设备时，内存复制操作被调用了两次。总共复制的字节数为 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操作，而是首先在两个流中排队内存复制操作，然后排队内核计算操作。这将确保内存复制和内核计算相互重叠。如果这两个操作花费相同的时间，我们可以实现两倍的速度提升。我们可以通过查看以下图表来更好地了解操作顺序：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/e068bb15-3dcf-4826-8563-cb0c96994f04.png

时间从上到下增加。我们可以看到，在同一时间段内执行了两个内存复制操作和内核执行操作，这将加速你的程序。我们还看到，由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函数释放。这是强制性的，否则您的系统可能会很快耗尽内存。程序输出如下所示：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/e2d48765-0845-4e40-8995-8f961ec3fc48.png

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

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/a545b18f-26fd-4176-9963-5af2a4752a5e.png

这是一个非常简单的情况，你可能看不到 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的值之间。第二个数组，用于存储直方图，被初始化为零。这两个数组被复制到设备内存中。内核将计算直方图并将其返回到主机。我们将在控制台上打印这个直方图。输出如下所示：

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/23a248c6-ceb8-47b0-ab44-bb8f91850aca.png

当我们尝试使用原子操作来测量这段代码的性能并与 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函数与前面的函数类似。

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

https://github.com/OpenDocCN/freelearn-ml-zh/raw/master/docs/hsn-gpu-acc-cv-ocv-cuda/img/2bfad7b5-3977-4a99-977d-d7ef4d5e290c.png

如果你测量前面程序的性能，它将击败没有共享内存的 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 加速概念。从下一章开始，我们将处理真实图像而不是矩阵。

问题

1. 为什么不使用 CPU 计时器来衡量内核函数的性能？

2. 尝试使用 Nvidia Visual Profiler 工具可视化上一章中实现的矩阵乘法代码的性能。

3. 给出程序中遇到的不同语义错误示例。

4. 内核函数中线程发散的缺点是什么？请用例子解释。

5. 使用cudahostAlloc函数在主机上分配内存的缺点是什么？

6. 证明以下陈述的正确性：CUDA 流中的操作顺序对于提高程序性能非常重要。

7. 为了在 CUDA 中实现良好的性能，对于 1024 x 1024 的图像，应该启动多少个块和线程？