作者：（美）包米克·维迪雅（Bhaumik Vaidya）
出版社：机械工业出版社
格式: AZW3, DOCX, EPUB, MOBI, PDF, TXT
基于GPU加速的计算机视觉编程：使用OpenCV和CUDA实时处理复杂图像数据试读：
前言
计算机视觉正在给许多行业带来革命性的变化，OpenCV是使用最广泛的计算机视觉工具之一，能够在多种编程语言中工作。目前，需要在计算机视觉中实时处理较大的图像，而单凭OpenCV是难以做到的。在这方面图形处理器（GPU）和CUDA可以提供帮助。因此，本书提供了一个将OpenCV与CUDA集成的实际应用的详细概述。本书首先解释了用CUDA对GPU编程，这对于从未使用过GPU的计算机视觉开发人员来说是必不可少的。然后通过一些实例说明了如何用GPU和CUDA对OpenCV进行加速。当要在现实生活场景中使用计算机视觉应用程序时，需要将其部署在嵌入式开发板上，本书涵盖了如何在NVIDIA Jetson TX1上部署OpenCV应用程序，这是非常流行的计算机视觉和深度学习应用程序。本书的最后一部分介绍了PyCUDA，结合Python使用OpenCV的计算机视觉开发人员会使用它。PyCUDA是一个Python库，它利用CUDA和GPU的功能来加速。本书为在C++或Python中使用OpenCV的开发人员提供了一个完整的指南，帮助他们通过亲身体验来加速计算机视觉应用程序。本书的读者对象
对于想学习如何利用GPU处理更复杂的图像数据的OpenCV开发人员，本书是必读指南。大多数计算机视觉工程师或开发人员在试图实时处理复杂的图像数据时都会遇到问题。这就需要使用GPU进行计算机视觉算法的加速，而这有助于人们开发能够实时处理复杂图像数据的算法。大多数人认为硬件加速只能通过FPGA和ASIC设计来实现，为此，他们需要Verilog或VHDL等硬件描述语言的知识。然而，只在CUDA出现之前情况才如此。CUDA利用了NVIDIA GPU的强大功能，可以使用支持CUDA的C++和Python等编程语言来加速算法。本书将通过开发实际应用程序来帮助这些开发人员了解这些概念并在嵌入式平台上部署计算机视觉应用程序，如NVIDIA Jetson TX1。本书的主要内容
第1章介绍了CUDA架构以及它如何重新定义GPU的并行处理能力，讨论了CUDA架构在实际场景中的应用，介绍了CUDA的开发环境，以及如何在所有操作系统上安装CUDA。
第2章教读者使用CUDA为GPU编写程序。从一个简单的Hello World程序开始，逐步用CUDA C构建复杂示例。该章还介绍了内核如何工作以及如何使用设备属性，并讨论了与CUDA编程相关的术语。
第3章向读者介绍了如何从CUDA程序中调用线程，多个线程如何相互通信，多个线程并行工作时如何同步，以及常量内存和纹理内存。
第4章包括CUDA流和CUDA事件等高级概念，描述了如何使用CUDA加速排序算法，并研究了使用CUDA加速简单图像处理功能。
第5章描述了在所有操作系统中安装支持CUDA的OpenCV库，解释了如何使用一个简单的程序来测试这个安装，比较了使用和不使用CUDA支持执行的图像处理程序的性能。
第6章教读者如何使用OpenCV开发基本的计算机视觉操作应用程序，如像素级的图像操作、滤波和形态学操作。
第7章介绍了使用OpenCV和CUDA加速一些实际计算机视觉应用程序的步骤，描述了用于对象检测的特征检测和描述算法。该章还介绍了基于Haar级联和视频分析技术的人脸检测加速，如用于对象跟踪的背景减法。
第8章介绍了Jetson TX1嵌入式平台以及如何使用它来加速和部署计算机视觉应用程序，还介绍了在Jetson TX1上使用JetPack安装文件安装OpenCV for Tegra的过程。
第9章包括在Jetson TX1上部署计算机视觉应用程序，介绍了如何构建不同的计算机视觉应用程序，以及如何将摄像机与Jetson TX1连接用于视频处理应用程序。
第10章介绍了PyCUDA，这是一个用于GPU加速的Python库。该章描述了在所有操作系统上的安装过程。
第11章教读者如何使用PyCUDA编写程序，其中详细描述了从主机到设备的数据传输和内核执行的概念，涵盖了如何在PyCUDA中使用数组和开发复杂的算法。
第12章介绍了使用PyCUDA的基本计算机视觉应用的开发和加速，描述了颜色空间转换操作、直方图计算和不同的算术操作作为计算机视觉应用的例子。充分利用本书
本书介绍的示例可以在Windows、Linux和macOS上运行，书中涵盖了所有的安装说明。读者最好对计算机视觉概念和编程语言（如C++和Python）有全面了解，最好用NVIDIA GPU硬件来执行书中介绍的示例。下载示例代码及彩色图像
本书的示例代码及所有截图和样图，可以从http://www.packtpub.com通过个人账号下载，也可以访问华章图书官网http://www.hzbook.com，通过注册并登录个人账号下载。
作者简介
Bhaumik Vaidya是一位经验丰富的计算机视觉工程师和导师，在OpenCV库尤其在计算机视觉问题解决方面做了大量工作。他是优秀硕士毕业生，目前正在攻读计算机视觉算法加速方面的博士学位，该算法使用OpenCV和基于GPU的深度学习库构建。他有教学背景，指导过许多计算机视觉和超大规模集成（VLSI）方面的项目。他之前在VLSI领域做过ASIC验证工程师，对硬件架构也有深入了解。他在著名期刊上发表了许多研究论文，还和博士导师共同获得了NVIDIA Jetson TX1嵌入式开发平台的研究资助。
审稿人简介
Vandana Shah于2001年获得电子学学士学位，随后获得了人力资源管理工商管理硕士和电子工程（超大规模集成电路方向）硕士学位，并提交了关于图像处理和脑肿瘤检测的深度学习领域的电子学博士论文。她擅长的领域是利用深度学习和嵌入式系统进行图像处理。她有超过13年的研究经验，并教育和指导着电子学和通信工程专业的本科生及研究生。曾在IEEE、Springer、Inderscience等出版的知名期刊发表论文。她还获得了政府资助，用于即将进行的核磁共振成像处理领域的研究。她一直致力于指导学生和研究人员，能够在软技能开发方面培训学生和教师。她多才多艺，除了拥有专业技术技能之外，还擅长印度舞蹈卡萨克。第1章 CUDA介绍及入门
本章向你简要介绍CUDA架构以及它是如何重新定义GPU的并行处理能力。应用软件如何使用CUDA架构？我们将演示一些实际的应用场景。本章希望成为使用通用GPU和CUDA加速的软件入门指南。本章描述了CUDA应用程序所使用的开发环境以及如何在各种操作系统上安装CUDA工具包。它涵盖了如何在Windows和Ubuntu上使用CUDA C开发基本代码。
本章将讨论以下主题：
·CUDA介绍
·CUDA应用
·CUDA开发环境
·在Windows、Linux和macOS上安装CUDA工具包
·使用CUDA C开发简单的代码1.1 技术要求
本章要求熟悉基本的C或C++编程语言。本章所有代码可以从GitHub链接https://github.com/PacktPublishing/Hands-On-GPU-Accelerated-Computer-Vision-with-OpenCV-and-CUDA下载。尽管代码只在Windows 10和Ubuntu 16.04上测试过，但可以在任何操作系统上执行。1.2 CUDA介绍
计算统一设备架构（Compute Unified Device Architecture，CUDA）是由英伟达（NVIDIA）开发的一套非常流行的并行计算平台和编程模型。它只支持NVIDIA GPU卡。OpenCL则用来为其他类型的GPU编写并行代码，比如AMD和英特尔，但它比CUDA更复杂。CUDA可以使用简单的编程API在图形处理单元（GPU）上创建大规模并行应用程序。
使用C和C++的软件开发人员可以通过使用CUDA C或C++来利用GPU的强大性能来加速他们的软件应用程序。用CUDA编写的程序类似于用简单的C或C++编写的程序，添加需要利用GPU并行性的关键字。CUDA允许程序员指定CUDA代码的哪个部分在CPU上执行，哪个部分在GPU上执行。
下一节将详细介绍并行计算的需求以及CUDA架构是如何利用GPU的强大性能。1.2.1 并行处理
近年来，消费者对手持设备的功能要求越来越高。因此，有必要将越来越多的晶体管封装在一个小的电路板上，既能快速工作，又能耗电最少。我们需要一个可以快速运行的处理器以较高的时钟速度、较小的体积和最小的功率执行多项任务。在过去的几十年中，晶体管的尺寸逐渐减小，这就可以让越来越多的晶体管封装在一个芯片上，也导致了时钟速度的不断提高。然而，这种情况已经发生了变化，最近几年时钟速度或多或少保持不变。那么，原因是什么呢？难道是晶体管不再变小了吗？答案是否定的。时钟速度恒定背后的主要原因是高功率损耗和高时钟速率。小的晶体管在小面积内封装和高速工作将耗散大功率，因此它是很难保持处理器的低温。开发中随着时钟速度逐渐饱和，我们需要一个新的计算模式来提高处理器性能。让我们通过一个真实生活中的小例子来理解这个概念。
假设你被告知要在很短的时间内挖一个很大的洞。你会有以下三种方法以及时完成这项工作：
·你可以挖得更快。
·你可以买一把更好的铲子。
·你可以雇佣更多的挖掘机，它们可以帮助你完成工作。
如果我们能在这个例子和一个计算模式之间找出关联，那么第一种选择类似于更快的时钟。第二种选择类似于拥有更多可以在每个时钟周期做更多工作的晶体管。但是，正如我们之前段落里讨论过的，功耗限制了这两个步骤。第三种选择是类似于拥有许多可以并行执行任务的更小更简单的处理器。GPU遵循这种计算模式。它不是一个可以执行复杂任务的更强大的处理器，而是有许多小而简单的且可以并行工作的处理器。下一节将解释GPU架构的细节。1.2.2 GPU架构和CUDA介绍
GeForce 256是英伟达于1999年开发的第一个GPU。最初只用在显示器上渲染高端图形。它们只用于像素计算。后来，人们意识到如果可以做像素计算，那么他们也可以做其他的数学计算。现在，GPU除了用于渲染图形图像外，还用于其他许多应用程序中。这些GPU被称为通用GPU（GPGPU）。
你可能会想到的下一个问题是CPU和GPU的硬件架构有什么不同，从而可以使得GPU能够进行并行计算？CPU具有复杂的控制硬件和较少的数据计算硬件。复杂的控制硬件在性能上提供了CPU的灵活性和一个简单的编程接口，但是就功耗而言，这是昂贵的。而另一方面，GPU具有简单的控制硬件和更多的数据计算硬件，使其具有并行计算的能力。这种结构使它更节能。缺点是它有一个更严格的编程模型。在GPU计算的早期，OpenGL和DirectX等图形API是与GPU交互的唯一方式。对于不熟悉OpenGL或DirectX的普通程序员来说，这是一项复杂的任务。这促成了CUDA编程架构的开发，它提供了一种与GPU交互的简单而高效的方式。关于CUDA架构的更多细节将在下一节中给出。
一般来说，任何硬件架构的性能都是根据延迟和吞吐量来度量的。延迟是完成给定任务所花费的时间，而吞吐量是在给定时间内完成任务的数量。这些概念并不矛盾。通常情况下，提高一个，另一个也会随之提高。在某种程度上，大多数硬件架构旨在提高延迟或吞吐量。例如，假设你在邮局排队。你的目标是在很短的时间内完成你的工作，所以你想要改进延迟，而坐在邮局窗口的员工想要在一天内看到越来越多的顾客。因此，员工的目标是提高吞吐量。在这种情况下，改进一个将导致另一个的改进，但是双方看待这个改进的方式是不同的。
同样，正常的串行CPU被设计为优化延迟，而GPU被设计为优化吞吐量。CPU被设计为在最短时间内执行所有指令，而GPU被设计为在给定时间内执行更多指令。GPU的这种设计理念使它们在图像处理和计算机视觉应用中非常有用，这也是本书的目的，因为我们不介意单个像素处理的延迟。我们想要的是在给定的时间内处理更多的像素，这可以在GPU上完成。
综上所述，如果我们想在相同的时钟速度和功率要求下提高计算性能，那么并行计算就是我们所需要的。GPU通过让许多简单的计算单元并行工作来提供这种能力。现在，为了与GPU交互，并利用其并行计算能力，我们需要一个由CUDA提供的简单的并行编程架构。1.2.3 CUDA架构
本节介绍在GPU架构中如何进行基本的硬件修改，以及使用CUDA开发程序的一般结构。我们暂时不讨论CUDA程序的语法，但是我们将讨论开发代码的步骤。本节还将介绍一些基本的术语，这些术语将贯穿全书。
CUDA架构包括几个专门为GPU通用计算而设计的特性，这在早期的架构中是不存在的。它包括一个unified shedder管道，它允许GPU芯片上的所有算术逻辑单元（ALU）被一个CUDA程序编组。ALU还被设计成符合IEEE浮点单精度和双精度标准，因此它可以用于通用应用程序。指令集也适合于一般用途的计算，而不是特定于像素计算。它还允许对内存的任意读写访问。这些特性使CUDA GPU架构在通用应用程序中非常有用。
所有的GPU都有许多被称为核心（Core）的并行处理单元。在硬件方面，这些核心被分为流处理器和流多处理器。GPU有这些流多处理器的网格。在软件方面，CUDA程序是作为一系列并行运行的多线程（Thread）来执行的。每个线程都在不同的核心上执行。可以将GPU看作多个块（Block）的组合，每个块可以执行多个线程。每个块绑定到GPU上的不同流多处理器。CUDA程序员不知道如何在块和流多处理器之间进行映射，但是调度器知道并完成映射。来自同一块的线程可以相互通信。GPU有一个分层的内存结构，处理一个块和多个块内线程之间的通信。这将在接下来的章节中详细讨论。
作为一名程序员，你会好奇CUDA中的编程模型是什么，以及代码将如何理解它是应该在CPU上执行还是在GPU上执行。本书假设我们有一个由CPU和GPU组成的计算平台。我们将CPU及其内存称为主机（Host），GPU及其内存称为设备（Device）。CUDA代码包含主机和设备的代码。主机代码由普通的C或C++编译器在CPU上编译，设备代码由GPU编译器在GPU上编译。主机代码通过所谓的内核调用调用设备代码。它将在设备上并行启动多个线程。在设备上启动多少线程是由程序员来决定的。
现在，你可能会问这个设备代码与普通C代码有何不同。答案是，它类似于正常的串行C代码。只是这段代码是在大量内核上并行执行的。然而，要使这段代码工作，它需要设备显存上的数据。因此，在启动线程之前，主机将数据从主机内存复制到设备显存。线程处理来自设备显存的数据，并将结果存储在设备显存中。最后，将这些数据复制回主机内存进行进一步处理。综上所述，CUDA C程序的开发步骤如下：
1）为主机和设备显存中的数据分配内存。
2）将数据从主机内存复制到设备显存。
3）通过指定并行度来启动内核。
4）所有线程完成后，将数据从设备显存复制回主机内存。
5）释放主机和设备上使用的所有内存。1.3 CUDA应用程序
CUDA在过去十年经历了前所未有的增长。它被广泛应用于各个领域的各种应用中。它改变了多个领域的研究。在本节中，我们将研究其中的一些领域，以及CUDA如何加速每个领域的增长：
·计算机视觉应用：计算机视觉和图像处理算法是计算密集型的。越来越多的摄像头在捕获高分辨率图像时，需要实时处理这些大图像。随着这些算法实现CUDA加速，图像分割、目标检测和分类等应用可以实现超过30帧/秒的实时帧率性能。CUDA和GPU允许对深度神经网络和其他深度学习算法进行更快的训练，这改变了计算机视觉的研究。英伟达正在开发多个硬件平台，如Jetson TK1、Jetson TX1和Jetson TX2，这些平台可以加速计算机视觉应用。英伟达drive平台也是为自动驾驶应用而设计的平台之一。
·医学成像：在医学成像领域，GPU和CUDA被广泛应用于磁共振成像和CT图像的重建和处理。这大大减少了这些图像的处理时间。现在，带有GPU的设备，可以借助一些库来使用CUDA加速处理这些图像。
·金融计算：所有金融公司都需要以更低的成本进行更好的数据分析，这将有助于做出明智的决策。它包括复杂的风险计算及初始和寿命裕度计算，这些都必须实时进行。GPU帮助金融公司在不增加太多间接成本的情况下，实时地做多种分析。
·生命科学、生物信息学和计算化学：模拟DNA基因、测序和蛋白质对接是需要大量计算资源的计算密集型任务。GPU有助于这种分析和模拟。GPU可以运行普通的分子动力学、量子化学和蛋白质对接应用程序，比普通CPU快5倍以上。
·天气研究和预报：与CPU相比，利用GPU和CUDA的几种天气预报应用、海洋建模技术和海啸预测技术可以进行更快的计算和模拟。
·电子设计自动化（EDA）：随着日益复杂的超大规模集成电路技术和半导体制造工艺的发展，使得EDA工具的性能在这一技术进步上落后。它导致了模拟不完整和功能bug的遗漏。因此，EDA行业一直在寻求更快的仿真解决方案。GPU和CUDA加速正在帮助这个行业加速计算密集型EDA模拟，包括功能模拟、placement和rooting，以及信号完整性和电磁学、SPICE电路模拟等。
·政府和国防：GPU和CUDA加速也被政府和军队广泛使用。航空航天、国防和情报工业正在利用CUDA加速将大量数据转化为可操作的信息。1.4 CUDA开发环境
要开始使用CUDA开发应用程序，你需要为它配置开发环境。为CUDA建立开发环境应具备以下先决条件：
·支持CUDA的GPU
·英伟达显卡驱动程序
·标准C编译器
·CUDA开发工具包
下面的几节将讨论如何检查第1个和第4个先决条件并安装它们。1.4.1 支持CUDA的GPU
如前所述，CUDA架构仅支持NVIDIA GPU。它不支持其他GPU，如AMD和英特尔。英伟达在过去十年中开发的几乎所有GPU都支持CUDA架构，可以用于开发和执行CUDA应用程序。可以在英伟达网站上找到支持CUDA的GPU的详细列表，网址为：https://developer.nvidia.com/cuda-gpus。如果你的GPU在列表里，那么你可以在你的PC上运行CUDA应用。
如果你不知道你的PC上是哪个GPU，可以通过以下步骤找到它：
在Windows下：
1）在开始菜单中，输入设备管理器，然后按Enter键。
2）在设备管理器中，单击显示适配器。在那里，你会找到你的NVIDIA GPU的名称。
在Linux下：
1）打开Terminal。
2）运行sudo lshw-C video。
这将列出有关显卡的信息，通常包括它的制造商和型号。
在macOS下：
1）到苹果菜单|关于这个Mac|更多信息。
2）在内容列表下选择图形/显示。在那里，你会找到你的NVIDIA GPU的名称。
如果你有一个支持CUDA的GPU，那么你可以继续下一步。1.4.2 CUDA开发工具包
CUDA需要一个GPU编译器来编译GPU代码。这个编译器附带一个CUDA开发工具包。如果你有一个最新驱动程序更新的NVIDIA GPU，并且为你的操作系统安装了一个标准的C编译器，那么你可以进入安装CUDA开发工具包的最后一步。下一节将讨论安装CUDA开发工具包的步骤。1.5 在所有操作系统上安装CUDA工具包
本节介绍了如何在所有支持平台安装CUDA工具包，以及如何验证是否安装成功。
安装CUDA时，可以选择下载在线安装器或离线本地安装器。前者需要手工下载的大小比较小，但是安装时需要连接互联网。后者一次性下载完成后虽然较大，但是安装时不需要连接互联网。可以从https://developer.nvidia.com/cuda-downloads下载适合Windows、Linux以及macOS的安装包。注意曾经有两种CUDA开发包（32位和64位），但现在NVIDIA已经放弃对32位版本的支持，因此你只能安装64位版本的。注意下面我们用CUDAx.x代表你实际下载到的CUDA工具包版本。
本书选择后者。1.5.1 Windows
本节介绍在Windows上安装CUDA的步骤，如下所示：
1）双击安装程序。它将要求你选择将提取临时安装文件的文件夹。选择你选择的文件夹。建议将此作为默认值。
2）然后，安装程序将检查系统兼容性。如果你的系统兼容，则可以按照屏幕提示安装CUDA。你可以选择快速安装（默认）和自定义安装。自定义安装允许选择要安装的CUDA功能。建议选择快速安装。
3）安装程序还将安装CUDA示例程序和CUDA Visual Studio集成。在运行之前，请确保已安装Visual Studio安装程序。1.5.2 Linux
本节介绍了如何在Linux发行版上安装CUDA开发包。Ubuntu是一种很流行的Linux发行版。具体的安装过程，将分别讨论使用NV提供的针对特定（Ubuntu）发行版的安装包和使用Ubuntu特定的apt-get命令这两种方式。
从前面的CUDA页面下载*.deb安装程序，然后按以下具体步骤安装：
1）打开终端并运行dpkg命令，该命令用于在基于Debian的系统中安装包：
2）使用以下命令安装CUDA公共GPG密钥：
3）使用以下命令更新apt repository缓存：
4）使用以下命令安装CUDA：
5）用下面的命令修改PATH环境变量，以包含CUDA安装路径的bin目录：如果你没有在默认位置安装CUDA，则用你的实际安装目录代替这里的例子。
6）通过这行命令设定LD_LIBRARY_PATH环境变量，来设定库搜索目录：
此外，你还可以通过第二种方式来安装CUDA开发包，也就是使用Ubuntu自带的apt-get。在命令行终端里输入如下命令即可：
7）nvcc将分别编译.cu文件中的Host和Device代码，前者是通过系统自带的GCC之类的Host代码编译器进行的，而后者则是通过CUDA C前端等一系列工具进行。你可以通过如下命令安装NSight Eclipse Edition（这也是NV的叫法），用作Linux下开发CUDA程序的图形化IDE环境。
安装后，你可以在用户Home目录下面，编译并执行~/NVIDIA_CUDA-x.x_Samples/下面的deviceQuery例子。如果你的CUDA开发包安装和配置正确的话，成功编译并运行后，你应当看到类似如图1-1所示的输出。图 1-11.5.3 Mac
本节介绍在macOS上安装CUDA的步骤。从CUDA网站下载*.dmg安装程序。下载安装程序后安装的步骤如下：
1）启动安装程序并按照屏幕提示完成安装。它将安装所有预制件、CUDA、工具包和CUDA示例。
2）需要使用以下命令设置环境变量：如果你没有在默认位置安装CUDA，则需要更改指向安装位置的路径。
3）运行脚本：cuda-install-samples-x.x.sh。它将安装具有写权限的CUDA示例。
4）完成之后，可以转到bin/x86_64/darwin/release并运行deviceQuery程序。如果CUDA工具包安装和配置正确，它将显示您的GPU的设备属性。1.6 一个基本的CUDA C程序
在本节中，我们将通过使用CUDA C编写一个非常基础的程序来学习CUDA编程。我们将从编写一个“Hello,CUDA!”开始，在CUDA C中编程并执行它。在详细介绍代码之前，有一件事你应该记得，主机代码是由标准C编译器编译的，设备代码是由NVIDIA GPU编译器来执行。NVIDIA工具将主机代码提供给标准的C编译器，例如Windows的Visual Studio和Ubuntu的GCC编译器，并使用macOS执行。同样需要注意的是，GPU编译器可以在没有任何设备代码的情况下运行CUDA代码。所有CUDA代码必须保存为*.cu扩展名。
下面就是Hello,CUDA!的代码
如果你仔细查看代码，它看起来将非常类似于简单地用C语言编写的Hello,CUDA!用于CPU执行的程序。这段代码的功能也类似。它只在终端或命令行上打印“Hello,CUDA!”。因此，你应该想到两个问题：这段代码有何不同？CUDA C在这段代码中扮演何种角色？这些问题的答案可以通过仔细查看代码来给出。它与用简单的C编写的代码相比，有两个主要区别：
·一个名为myfirstkernel的空函数，前缀为__global__
·使用<<1,1>>>调用myfirstkernel函数
__global__是CUDA C在标准C中添加的一个限定符，它告诉编译器在这个限定符后面的函数定义应该在设备上而不是在主机上运行。在前面的代码中，myfirstkernel将运行在设备上而不是主机上，但是，在这段代码中，它是空的。
那么，main函数将在哪里运行？NVCC编译器将把这个函数提供给C编译器，因为它没有被global关键字修饰，因此main函数将在主机上运行。
代码中的第二个不同之处在于对空的myfirstkernel函数的调用带有一些尖括号和数值。这是一个CUDA C技巧：从主机代码调用设备代码。它被称为内核调用。内核调用的细节将在后面的章节中解释。尖括号内的值表示我们希望在运行时从主机传递给设备的参数。基本上，它表示块的数量和将在设备上并行运行的线程数。因此，在这段代码中，<<<1,1>>>表示myfirstkernel将运行在设备上的一个块和一个线程或块上。虽然这不是对设备资源的最佳使用，但是理解在主机上执行的代码和在设备上执行的代码之间的区别是一个很好的起点。
让我们再来重温和修改“Hello,CUDA!”代码，myfirstkernel函数将运行在一个只有一个块和一个线程或块的设备上。它将通过一个称为内核启动的方法从main函数内部的主机代码启动。
在编写代码之后，你将如何执行此代码并查看输出？下一节将描述在Windows和Ubuntu上编写和执行Hello CUDA!代码的步骤！1.6.1 在Windows上创建CUDA C程序的步骤
本节描述使用Visual Studio在Windows上创建和执行基本CUDA C程序的步骤。步骤如下：
1）打开Microsoft Visual Studio。
2）进入File|New|Project。
3）依次选择NVIDIA|CUDA 9.0|CUDA 9.0 Runtime。
4）为项目自定义名称，然后单击OK按钮。
5）它将创建一个带有kernel.cu示例文件的项目。现在双击打开这个文件。
6）从文件中删除现有代码，写入前面编写的那段代码。
7）从生成（Build）选项卡中选择生成（build）进行编译，并按快捷键Ctrl+F5调试代码。如果一切正常，你会看到Hello,CUDA!显示在命令行上，如图1-2所示。图 1-21.6.2 在Ubuntu上创建CUDA C程序的步骤
本节描述使用Nsight Eclipse插件在Ubuntu上创建和执行基本CUDA C程序的步骤。步骤如下：
1）打开终端并输入nsight来打开Nsight。
2）依次选择File|New|CUDA C/C++Projects。
3）为项目自定义名称，然后单击OK按钮。
4）它将创建一个带有示例文件的项目。现在双击打开这个文件。
5）从文件中删除现有代码，写入前面编写的那段代码。
6）按下play按钮运行代码。如果一切正常，你会看到Hello,CUDA!显示在终端，如图1-3所示。图 1-31.7 总结
在这一章中，我介绍了CUDA，并简要介绍了并行计算的重要性。我们还详细讨论了CUDA和GPU在各个领域的应用。本章描述了在PC上执行CUDA应用程序所需的硬件和软件设置。我们给出了在本地PC上安装CUDA的详细步骤。
1.6节通过开发一个简单的程序并在Windows和Ubuntu上执行，给出了CUDA C中的应用程序开发的入门指南。
在下一章中，我们将基于CUDA C中的编程知识，通过几个实际示例介绍使用CUDA C的并行计算，以展示它如何比普通编程更快。还将介绍线程和块的概念，以及如何在多线程和块之间执行同步。1.8 测验题
1.解释三种提高计算硬件性能的方法。使用哪种方法开发GPU？
2.真假判断：改进延迟将提高吞吐量。
3.填空：CPU被设计用来改进____，GPU被设计用来改进___。
4.举个例子，从一个地方到240公里以外的另一个地方。你可以开一辆能容纳5人的车，时速60公里，或者开一辆能容纳40人的公交车，时速40公里。哪个选项将提供更好的延迟，哪个选项将提供更好的吞吐量？
5.解释GPU和CUDA在计算机视觉应用中特别有用的原因。
6.真假判断：CUDA编译器不能在没有设备代码的情况下编译代码。
7.在本章讨论的“Hello,CUDA!”例子中，printf语句是由主机执行还是由设备执行的？第2章 使用CUDA C进行并行编程
在上一章中，我们看到了安装CUDA并使用它编写程序是多么容易。尽管这个示例并不令人印象深刻，但它证明了使用CUDA是非常容易的。在本章中，我们将以这个概念为基础，教你如何使用CUDA为GPU编写高级程序。我们从变量加法程序开始，然后逐步构建CUDA C中的复杂向量操作示例，我们会介绍内核如何工作以及如何在CUDA程序中使用设备属性。本章还会讨论在CUDA程序中向量是如何运算的，以及与CPU处理相比，CUDA如何能加速向量运算。除此之外，我们还会介绍与CUDA编程相关的术语。
本章将讨论以下主题：
·内核调用的概念
·在CUDA中创建内核函数并向其传递参数
·配置CUDA程序的内核参数和内存分配
·CUDA程序中的线程执行
·在CUDA程序访问GPU设备属性
·在CUDA程序中处理向量
·并行通信模型2.1 技术要求
本章要求熟悉基本的C或C++编程语言，特别是动态内存分配。本章所有代码可以从GitHub链接https://github.com/PacktPublishing/Hands-On-GPU-Accelerated-Computer-Vision-with-OpenCV-and-CUDA下载。尽管代码只在Windows 10和Ubuntu 16.04上测试过，但可以在任何操作系统上执行。2.2 CUDA程序结构
前面我们看到了一个非常简单的“Hello,CUDA!”程序，其中展示了一些与CUDA程序相关的重要概念。CUDA程序是在主机或GPU设备上执行的函数的组合。不显示并行性的函数在CPU上执行，显示数据并行性的函数在GPU上执行。GPU编译器在编译期间隔离这些函数。如前一章所示，在设备上执行的函数是使用__global__关键字定义的，由NVCC编译器编译，而普通的C主机代码是由C编译器编译的。CUDA代码基本上与ANSI C代码相同，只是添加了一些开发数据并行性所需的关键字。
因此，在本节中，我们用一个简单的双变量加法程序来解释与CUDA编程相关的重要概念，如内核调用、从主机到设备传递参数到内核函数、内核参数的配置、利用数据并行性需要的CUDA API，以及发生在主机和设备上的内存分配。2.2.1 CUDA C中的双变量加法程序
在第1章里，我们演示了一个简单的“Hello,CUDA!”代码，里面的设备函数是空的，这无关紧要。本节介绍一个简单的加法程序，它在设备上执行两个变量的加法。虽然它没有利用设备的任何数据并行性，但它对于演示CUDA C的重要编程概念非常有用。首先，我们将看到如何编写一个将两个变量相加的内核（kernel）函数。
内核函数代码如下：
gpuAdd函数与ANSI C中的一个普通add函数非常相似。它以两个整数变量d_a和d_b作为输入，并将加法存储在第三个整数指针d_c所指示的内存位置。设备函数的返回值为void，因为它将结果存储在设备指针指向的内存位置中，而不显式地返回任何值。现在我们将看到如何为这段代码编写main函数。main函数代码如下：
在main函数中，前两行定义主机和设备的变量。第三行使用cudaMalloc函数在设备上分配d_c变量的内存。cudaMalloc函数类似于C中的malloc函数。在main函数的第四行中，调用gpuAdd，其中1和4是两个输入变量，d_c是一个作为输出指针变量的设备显存指针。gpuAdd函数的独特语法（也称为内核调用）将在下一节中解释。如果gpuAdd的结果需要在主机上使用，那么它必须从设备的内存复制到主机的内存中，这是由cudaMemcpy函数完成的。然后，使用printf函数打印这个结果。倒数第二行使用cudaFree函数释放设备上使用的内存。从程序中释放设备上使用的所有内存是非常重要的，否则，你可能在某个时候耗尽内存。以//开头的行是使代码可读性更高的注释，编译器会忽略这些行。
双变量加法程序有两个函数：main和gpuAdd。如你所见，gpuAdd是通过使用__global__关键字定义的，因此它用于在设备上执行，而main函数将在主机上执行。这个程序将设备上的两个变量相加，并在命令行上打印输出，如图2-1所示。
我们将在本书中使用一个约定，主机变量将以h_为前缀，设备变量将以d_为前缀。这不是强制性的，这样做只是为了让读者能够轻松地理解概念，而不会混淆主机和设备。图 2-1
所有CUDA API（如cudaMalloc、cudaMemcpy和cudaFree）以及其他重要的CUDA编程概念（如内核调用、向内核传递参数以及内存分配问题）将在后面的部分中讨论。2.2.2 内核调用
使用ANSI C关键字和CUDA扩展关键字编写的设备代码称为内核。它是主机代码（Host Code）通过内核调用的方式来启动的。简单地说，内核调用的含义是我们从主机代码启动设备代码。内核调用通常会生成大量的块（Block）和线程（Thread）来在GPU上并行地处理数据。内核代码非常类似于普通的C函数，只是这段代码是由多个线程并行执行的。内核启动的语法比较特殊，如下所示：
它以我们想要启动的内核的名称开始。你应该确保这个内核是使用__global__关键字定义的。然后，它具有<<<>>>内核启动配置，该配置包含内核的配置参数。它可以包含三个用逗号分隔的参数。第一个参数表示希望执行的块数，第二个参数表示每个块将具有的线程数。因此，内核启动所启动的线程总数就是这两个数字的乘积。第三个参数是可选的，它指定内核使用的共享内存的大小。在变量相加程序中，内核启动语法如下：
在这里，gpuAdd是我们想要启动的内核的名称，<<<1,1>>>表示我们想用每个块一个线程启动一个块，这意味着我们只启动一个线程。圆括号中的三个参数是传递给内核的参数。这里，我们传递了两个常数，1和4。第三个参数是指向d_c设备显存的指针。它指向设备显存中的位置，内核将在那里存储相加后的结果。程序员必须记住的一件事是，作为参数传递给内核的指针应该仅指向设备显存。如果它指向主机内存，会导致程序崩溃。内核执行完成后，设备指针指向的结果可以复制回主机内存，以供进一步使用。只启动一个线程在设备上执行不是设备资源的最佳使用。假设你想并行启动多个线程，在内核调用的语法中需要做哪些修改？这将在下一节讨论，并称为“配置内核参数”。2.2.3 配置内核参数
为了在设备上并行启动多个线程，我们必须在内核调用中配置参数，内核调用是在内核启动配置中编写的。它们指定了Grid中块的数量，和每个块中线程的数量。我们可以并行启动很多个块，而每个块内又有很多个线程。通常，每个块有512或1024个线程。每个块在流多处理器上运行，一个块中的线程可以通过共享内存（Shared Memory）彼此通信。程序员无法选定哪个流多处理器将执行特定的块，也无法选定块和线程以何种顺序执行。
假设要并行启动500个线程，你可以对前面解释的内核启动语法进行哪些修改？一种选择是通过以下语法启动一个包含500个线程的块：
我们还可以启动一个线程的500个块，或者两个线程，每个线程250个块。因此，你必须修改内核启动配置里的值。程序员必须注意，每个块的线程数量不能超过GPU设备所支持的最大限制。在本书中，我们的目标是计算机视觉应用程序，需要处理二维和三维图像。在这里，如果块和线程不是一维的，而是多维的，那就可以更好地进行处理和可视化。
GPU支持三维网格块和三维线程块。它有以下语法：
在这里，N、N和N分别表示网格中沿x，y和z轴方向的块数。bxbybz同样，N、N和N分别表示一个块中沿x，y和z轴方向的线程数。txtytz如果没有指定y和z的维数，默认情况下它们被取为1。例如，要处理一个图像，你可以启动一个16×16的块网格，所有的块都包含16×16个线程。语法如下：
总之，在启动内核时，块数量和线程数量的配置非常重要。根据我们正在开发的应用程序和GPU资源的不同，应该谨慎地选择。下一节将解释在常规ANSI C函数上添加的一些重要CUDA函数。2.2.4 CUDA API函数
在变量加法程序中，我们遇到了一些常规C或C++程序员不熟悉的函数或关键字。这些关键字和函数包括__global__、cudaMalloc、cudaMemcpy和cudaFree。因此，在本节中，我们将逐一详细介绍这些函数。
·__global__：它与__device__和__host__一起是三个限定符关键字。这个关键字表示一个函数被声明为一个设备函数，当从主机调用时将在设备上执行。应该记住，这个函数只能从主机调用。如果要在设备上执行函数并从设备函数调用函数，则必须使用__device__关键字。__host__关键字用于定义只能从其他主机函数调用的主机函数。这类似于普通的C函数。默认情况下，程序中的所有函数都是主机函数。__host__和__device__都可以同时用于定义任何类型函数。它生成同一个函数的两个副本。一个将在主机上执行，另一个将在设备上执行。
·cudaMalloc：它类似于C中用于动态内存分配的Malloc函数。此函数用于在设备上分配特定大小的内存块。举例说明cudaMalloc的语法如下：
如上面的示例代码所示，它分配了一个大小等于一个整数变量大小的内存块，并返回指向该内存位置的指针d_c。
·cudaMemcpy：这个函数类似于C中的Memcpy函数，用于将一个内存区域复制到主机或设备上的其他区域。它的语法如下：
这个函数有四个参数。第一个参数是目标指针，第二个参数是源指针，它们分别指向主机内存或设备显存位置。第三个参数表示数据复制的大小，最后一个参数表示数据复制的方向：可以从主机到设备，设备到设备，主机到主机，或设备到主机。但是要小心，前两个指针参数必须和这里的复制方向参数是一致的。如示例所示，通过指定设备指针d_c作为源指针，主机指针h_c作为目标，我们将一个整数变量，从设备显存复制到了主机内存上。
·cudaFree：类似于C中的free函数，cudaFree的语法如下：
它释放了d_ptr指向的内存空间。在示例代码中，它释放了d_c指向位置的内存。请确保分配了d_c内存，使用了cudaMalloc，再使用cudaFree释放它。
CUDA除了现有的ANSI C函数之外，还有许多其他的关键字和函数。我们会经常使用这三个函数，因此本节对它们进行了讨论。要了解更多细节，你可以访问CUDA编程指南。2.2.5 将参数传递给CUDA函数
变量加法程序的gpuAdd内核函数与普通的C函数非常相似。因此，与普通的C函数一样，内核函数也可以按值或引用传递参数。因此，在本节中，我们将看到传递CUDA内核参数的两个方法。1.按值传递参数
回忆一下，在gpuAdd程序中，调用内核的语法如下：
另一方面，定义中的gpuAdd函数原型如下：
因此，你可以看到我们在调用内核时传递了d_a和d_b的值。首先，参数1会被复制到d_a，然后参数4会在调用内核时被复制到d_b。相加后的结果将存储在设备显存中d_c指向的地址。与其直接将值1和4作为输入传递给内核，我们还可以这样写：
在这里，a和b是可以包含任何整数值的整数变量。不建议按值传递参数，因为这会在程序中造成不必要的混乱和复杂性。最好是修改后的通过引用传递参数。
试读结束[说明：试读内容隐藏了图片]