CUDA-编程学习手册（全）

CUDA 编程学习手册（全）

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

译者：飞龙

协议：CC BY-NC-SA 4.0

前言

不要在第一次胜利后就休息。因为如果你在第二次失败，会有更多的人说你的第一次胜利只是运气。

A. P. J. Abdul Kalam

传统上，计算需求与中央处理单元（CPU）相关联，CPU 已经从单核发展到现在的多核。每一代新的 CPU 都提供了更多的性能，但科学和高性能计算社区每年都要求更多的性能，导致应用程序需求与硬件/软件堆栈提供的计算之间存在差距。与此同时，传统上用于视频图形的新架构也进入了科学领域。图形处理单元（GPU）——基本上是用于加速计算机图形的并行计算处理器——在 2007 年 CUDA 推出时进入了 HPC 领域。CUDA 成为了使用 GPU 进行通用计算的事实标准；即非图形应用程序。

自 CUDA 诞生以来，已经发布了许多版本，现在 CUDA 的版本为 10.x。每个版本都提供支持新硬件架构的新功能。本书旨在帮助您学习 GPU 并行编程，并指导您在现代应用中的应用。在它的帮助下，您将能够发现现代 GPU 架构的 CUDA 编程方法。本书不仅将指导您了解 GPU 功能、工具和 API，还将帮助您了解如何使用示例并行编程算法来分析性能。本书将确保您获得丰富的优化经验和对 CUDA 编程平台的洞察，包括各种库、开放加速器（OpenACC）和其他语言。随着您的进步，您将发现如何在一个盒子或多个盒子中利用多个 GPU 生成额外的计算能力。最后，您将探索 CUDA 如何加速深度学习算法，包括卷积神经网络（CNN）和循环神经网络（RNN）。

本书旨在成为任何新手或初学者开发者的入门点。但到最后，您将能够为不同领域编写优化的 CUDA 代码，包括人工智能。

如果您符合以下情况，这本书将是一个有用的资源：

您是 HPC 或并行计算的新手
您有代码，并希望通过将并行计算应用于 GPU 来提高其性能
您是深度学习专家，想利用 GPU 加速深度学习算法，如 CNN 和 RNN
您想学习优化代码和分析 GPU 应用性能的技巧和窍门，并发现优化策略
您想了解最新的 GPU 功能，以及高效的、分布式的多 GPU 编程。

如果您觉得自己属于以上任何一类，请加入我们一起踏上这段旅程。

这本书适合谁

这本初学者级别的书适用于希望深入研究并行计算、成为高性能计算社区的一部分并构建现代应用程序的程序员。假定具有基本的 C 和 C++编程经验。对于深度学习爱好者，本书涵盖了 Python InterOps、DL 库以及性能估计的实际示例。

为了充分利用这本书

本书适用于完全初学者和刚开始学习并行计算的人。除了计算机体系结构的基础知识外，不需要任何特定的知识，假定具有 C/C++编程经验。对于深度学习爱好者，在[第十章]（d0e9e8ff-bc17-4031-bb0e-1cfd310aff6f.xhtml）中，还提供了基于 Python 的示例代码，因此该章节需要一些 Python 知识。

本书的代码主要是在 Linux 环境中开发和测试的。因此，熟悉 Linux 环境是有帮助的。任何最新的 Linux 版本，如 CentOS 或 Ubuntu，都可以。代码可以使用 makefile 或命令行进行编译。本书主要使用免费软件堆栈，因此无需购买任何软件许可证。本书中将始终使用的两个关键软件是 CUDA 工具包和 PGI 社区版。

由于本书主要涵盖了利用 CUDA 10.x 的最新 GPU 功能，为了充分利用所有培训材料，最新的 GPU 架构（Pascal 及更高版本）将是有益的。虽然并非所有章节都需要最新的 GPU，但拥有最新的 GPU 将有助于您重现本书中实现的结果。每一章都有一个关于首选或必备 GPU 架构的部分，位于技术要求部分。

下载示例代码文件

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

您可以按照以下步骤下载代码文件：

在www.packt.com上登录或注册。
选择支持选项卡。
点击“代码下载”。
在搜索框中输入书名，然后按照屏幕上的说明操作。

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

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

该书的代码包也托管在 GitHub 上，网址为github.com/PacktPublishing/Learn-CUDA-Programming。如果代码有更新，将在现有的 GitHub 存储库上进行更新。

我们还提供了来自我们丰富书籍和视频目录的其他代码包，可在github.com/PacktPublishing/上找到。快来看看吧！

下载彩色图像

我们还提供了一个 PDF 文件，其中包含本书中使用的屏幕截图/图表的彩色图像。您可以在这里下载：static.packt-cdn.com/downloads/9781788996242_ColorImages.pdf。

使用的约定

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

CodeInText：表示文本中的代码词、数据库表名、文件夹名、文件名、文件扩展名、路径名、虚拟 URL、用户输入和 Twitter 句柄。以下是一个例子：“请注意，cudaMemcpy有一个异步的替代方案。”

代码块设置如下：

#include<stdio.h>
#include<stdlib.h>

__global__ void print_from_gpu(void) {
    printf("Hello World! from thread [%d,%d] \
        From device\n", threadIdx.x,blockIdx.x);
}

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

int main(void) {
    printf("Hello World from host!\n");
    print_from_gpu<<<1,1>>>();
    cudaDeviceSynchronize();
    return 0;
}

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

$ nvcc -o hello_world hello_world.cu

粗体：表示新术语、重要单词或屏幕上看到的单词。例如，菜单中的单词或对话框中的单词会以这种形式出现在文本中。以下是一个例子：“对于 Windows 用户，在 VS 项目属性对话框中，您可以在 CUDA C/C++ | Device | Code Generation 中指定您的 GPU 的计算能力。”

警告或重要说明如下。

提示和技巧如下。

第一章：CUDA 编程简介

自 2007 年首次发布以来，统一计算设备架构（CUDA）已经成长为使用图形计算单元（GPU）进行通用计算的事实标准，即非图形应用程序。那么，CUDA 到底是什么？有人可能会问以下问题：

这是一种编程语言吗？
这是一个编译器吗？
这是一种新的计算范式吗？

在本章中，我们将揭开有关 GPU 和 CUDA 的一些神秘之谜。本章通过提供对高性能计算（HPC）历史的简化视图，并用摩尔定律和丹纳德缩放等法则加以证实，为异构计算奠定基础，这些法则一直在推动半导体行业，因此也推动了处理器架构本身。您还将了解 CUDA 编程模型，并了解 CPU 和 GPU 架构之间的根本区别。通过本章的学习，您将能够使用 C 语言中的 CUDA 编程构造编写和理解Hello World!程序。

尽管本章主要使用 C 语言来演示 CUDA 构造，但我们将在其他章节中涵盖其他编程语言，如 Python、Fortran 和 OpenACC。

本章将涵盖以下主题：

高性能计算的历史
来自 CUDA 的 Hello World
使用 CUDA 进行矢量加法
CUDA 的错误报告
CUDA 中的数据类型支持

高性能计算的历史

高性能计算一直在不断突破极限，以实现科学发现。处理器架构和设计的根本转变有助于跨越 FLOP 障碍，从百万浮点运算（MFLOPs）开始，现在能够在一秒内进行 PetaFLOP 计算。

每秒浮点运算（FLOPs）是衡量任何计算处理器理论峰值的基本单位。MegaFLOP 代表 FLOPS 的 10 的 6 次方。PetaFLOP 代表 FLOPS 的 10 的 15 次方。

指令级并行性（ILP）是一个概念，其中独立于代码的指令可以同时执行。为了使指令并行执行，它们需要彼此独立。所有现代 CPU 架构（甚至 GPU 架构）都提供了五到 15 个以上的阶段，以实现更快的时钟频率：

Instr 1: add = inp1 + inp2

Instr 2: mult = inp1 * inp2

Instr 3: final_result = mult / add

用于计算mult和add变量的操作不相互依赖，因此可以在计算final_result时同时计算，而final_result依赖于Instr 1和Instr 2操作的结果。因此，在计算add和mult之前无法计算它。

当我们从技术变革的角度看高性能计算的历史，这些变革导致了新处理器设计的根本转变，以及对科学界的影响，有三个主要的变革可以被称为时代：

时代 1：超级计算机的历史可以追溯到 CRAY-1，它基本上是一个提供峰值 160 MegaFLOP/MFLOP 计算能力的单一矢量 CPU 架构。
时代 2：通过从单核设计转向 CRAY-2 的多核设计，跨越了 MegaFLOP 障碍，CRAY-2 是一个 4 核矢量 CPU，提供了 2 GigaFLOPs 的峰值性能。
时代 3：跨越 GigaFLOP 计算性能是一个根本性的转变，需要计算节点相互协作，并通过网络进行通信，以提供更高的性能。 Cray T3D 是第一批提供 1 TeraFLOP 计算性能的机器之一。网络是 3D Torus，提供 300 MB/s 的带宽。这是标准微处理器周围丰富shell的第一个重要实现。

此后，将近 20 年没有根本性的创新。技术创新主要集中在三个架构创新上：

从 8 位到 16 位再到 32 位，现在是 64 位指令集
增加 ILP
增加核心数量

这得到了时钟频率的增加，目前为 4 GHz。由于驱动半导体行业的基本定律，这是可能的。

摩尔定律：这个定律观察到密集集成电路中的晶体管数量每两年翻一番。

摩尔的预测在几十年来一直准确无误。摩尔定律是对历史趋势的观察和预测。

Dennard 缩放：这是一个使摩尔定律保持活力的缩放定律。Dennard 观察到晶体管尺寸和功率密度之间的关系，并用以下公式总结了这一观察：

P = QfCV² + V I[leakage]

在这个方程中，Q是晶体管数量，f是操作频率，C是电容，V是操作电压，I[leakage]是泄漏电流。

Dennard 缩放和摩尔定律彼此相关，因为可以推断出，减小晶体管的尺寸可以在成本效益方面导致芯片上的晶体管数量越来越多。

根据 Dennard 缩放规则，对于给定尺寸的芯片，多个处理器世代的总芯片功率保持不变。晶体管数量翻倍，而尺寸不断缩小（1/S速率），并且每两年以 40%的速度增加频率。当特征尺寸达到 65 纳米以下时，这种情况停止了，因为泄漏电流呈指数增长，这些规则不再能够持续。为了减少泄漏电流的影响，新的创新被强制执行在开关过程中。然而，这些突破仍然不足以恢复电压的缩放。电压在许多处理器设计中保持在 1V 恒定。不再可能保持功率包络恒定。这也被称为 Powerwall。

Dennard 缩放从 1977 年一直持续到 1997 年，然后开始衰退。因此，从 2007 年到 2017 年，处理器从 45 纳米变为 16 纳米，但导致每芯片能耗增加了三倍。

同时，管线阶段从五个阶段发展到了最新架构的 15+个阶段。为了保持指令管线的充分，使用了先进的技术，比如推测。推测单元涉及预测程序的行为，比如预测分支和内存地址。如果预测准确，它可以继续；否则，它会撤销已经完成的工作并重新开始。深层管线阶段和传统软件的编写方式导致了未使用的晶体管和浪费的时钟周期，这意味着应用性能没有改善。

然后出现了 GPU，最初主要用于图形处理。研究人员马克·哈里斯首次利用 GPU 进行了非图形任务，并创造了新术语使用 GPU 进行通用计算（GPGPU）。GPU 在某些数据并行类任务方面被证明是有效的。毫不奇怪，许多 HPC 应用程序中的大部分计算密集型任务在性质上都是数据并行的。它们主要是矩阵乘法，这在基本线性代数规范（BLAS）中是常规且广泛使用的。

用户在适应和使用 GPU 时唯一的问题是他们必须了解图形管线以利用 GPU。提供给 GPU 上任何计算工作的唯一接口围绕着着色器执行。需要提供一个更通用的接口，让在 HPC 社区工作的开发人员熟悉。这个问题在 2007 年引入 CUDA 时得到解决。

虽然 GPU 架构也受到相同的定律约束（摩尔定律和 Dennard 缩放），但处理器的设计采用了不同的方法，为不同的用途专门分配晶体管，并实现了比传统的同质架构更高的性能。

以下图表显示了计算机体系结构从顺序处理到分布式内存的演变及其对编程模型的影响：

随着 GPU 被添加到现有服务器上，应用程序在两种处理器（CPU 和 GPU）上运行，引入了异构的概念。这是我们将在下一节介绍的内容。

异构计算

围绕 GPU 的一个常见误解是它是 CPU 的替代品。GPU 用于加速代码中并行的部分。加速器是一个常用术语，用于描述 GPU，因为它们通过更快地运行代码的并行部分来加速应用程序，而 CPU 运行另一部分代码，即延迟绑定的部分。因此，高效的 CPU 与高吞吐量的 GPU 相结合，可以提高应用程序的性能。

以下图表代表了在多种处理器类型上运行的应用程序：

这个概念可以很好地用阿姆达尔定律来定义。阿姆达尔定律用于定义当应用程序的一部分被并行化时可以实现的最大加速。为了演示这一点，前面的图表显示了代码的两个部分。一个部分是延迟绑定的，而另一个是吞吐量绑定的。我们将在下一节中介绍这两个术语，区分 CPU 和 GPU 体系结构。

关键点是，CPU 对于某些延迟绑定的代码部分很好，而 GPU 擅长并行运行代码的单指令多数据（SIMD）部分。如果在优化后只有其中一个，即 CPU 代码或 GPU 代码，运行速度更快，这不一定会导致整体应用程序的速度提升。需要的是，当两个处理器都得到最佳利用时，性能方面才能获得最大的好处。这种从处理器上卸载某些类型的操作到 GPU 的方法被称为异构计算。

以下图表描述了所有应用程序具有的两种部分，即延迟绑定和吞吐量绑定：

在这里，使用阿姆达尔定律演示了改进两个部分的重要性。

编程范式

计算机体系结构的分类是使用弗林分类法进行的，描述了四类体系结构。弗林的分类之一 SIMD 用于描述 GPU 体系结构。然而，两者之间存在微妙的差异。SIMD 用于描述同一指令并行应用于多个数据点的体系结构。这种描述适用于具有矢量化能力的处理器。相比之下，在单指令多线程（SIMT）中，不是单个线程发出指令，而是多个线程向不同的数据发出相同的指令。与 SIMD 相比，GPU 体系结构更适合 SIMT 类别。

让我们看一个例子，将两个数组相加并将数据存储在第三个数组中。这个操作的数据集包括数组A、B和C。用于加法的相同操作被用于数组的每个元素：

Cx = Ax + Bx

显然，每个任务都是独立的，但所有线程都在应用相同的操作。

以下截图显示了矢量加法，展示了这种范例的一个例子：

低延迟与高吞吐量

正如我们在前一节中提到的，CPU 架构被优化用于低延迟访问，而 GPU 架构被优化用于数据并行吞吐量计算。如下截图所示，与 GPU 相比，CPU 架构具有大量缓存并且具有许多类型。我们越高，即从 L3 到 L1，缓存的数量就越少，但延迟就越低。CPU 架构旨在实现对缓存数据集的低延迟访问。大量晶体管用于实现推测执行和乱序执行。由于 CPU 以非常高的时钟速度运行，因此有必要通过频繁地将使用的数据存储在缓存中并预测下一条要执行的指令来隐藏获取数据的延迟。可以最佳地利用 CPU 缓存的应用程序可以探索这种时间局部性。此外，可以利用填充指令管线的应用程序，例如代码中没有if和else语句的应用程序，通过隐藏获取指令的延迟来受益。因此，CPU 架构是一种减少延迟的架构。

以下截图显示了 CPU 和 GPU 架构如何为不同的内存和计算单元分配芯片芯片区域。GPU 使用大量晶体管进行计算 ALUs，而 CPU 使用它来减少延迟。

另一方面，GPU 架构被称为“减少延迟”或“高吞吐量架构”。GPU 架构通过来自其他线程的计算来隐藏延迟。当一个线程在等待数据可用进行计算时，其他线程可以开始执行，因此不会浪费任何时钟周期。如果您熟悉 CUDA，那么您可能已经了解到 warp 的概念。我们将在接下来的章节中介绍 warp 的概念。（在 CUDA 中，执行单元是 warp 而不是线程。因此，上下文切换发生在 warp 而不是线程之间）。

有些人可能已经在想为什么我们不能在 CPU 中创建这些线程并做同样的事情来隐藏延迟。原因是 GPU 有大量的寄存器，并且所有线程上下文切换信息已经存在于其中。这是最快的内存。然而，在 CPU 中，寄存器集是有限的，因此线程相关的信息通常存储在较低的内存层次结构中，比如缓存。例如，Volta 包含 20MB 的寄存器存储。因此，与 GPU 相比，CPU 中线程之间的上下文切换时间要长得多。

现在，让我们来看看在 GPU 编程方面的不同方法。

GPU 的编程方法

现在，让我们回到我们最初的问题，即 CUDA 是什么？CUDA 是由 NVIDIA 开发的并行计算平台和编程模型架构，它将 GPU 上的通用计算作为一流能力进行暴露。与任何其他处理器一样，GPU 架构可以使用各种方法进行编码。提供快速加速的最简单方法是利用现有库。另外，开发人员可以选择使用 OpenACC 指令以获得快速加速结果和可移植性。另一种选择是选择通过使用 C、C++、Fortran、Python 等语言构造来深入研究 CUDA，以获得最高的性能和灵活性。我们将在接下来的章节中详细介绍所有这些方法。

以下截图表示了我们可以进行 GPU 编程的各种方式。

在本节中，我们为您提供了处理器和高性能计算随时间演变的视角。我们为您提供了异构编程模型对于从应用程序中获得最佳性能的关键性概述，以及 GPU 编程的方法。在下一节中，我们将开始在 GPU 上编写一个 Hello World 程序。

技术要求

本章需要一台安装了现代 NVIDIA GPU（Pascal 架构及以上）的 Linux/Windows PC，以及所有必要的 GPU 驱动程序和安装了 CUDA Toolkit（10.0 及以上版本）。如果您不确定您的 GPU 架构，请访问 NVIDIA 的 GPU 网站（developer.nvidia.com/cuda-gpus）并确认您的 GPU 架构。本章的代码也可以在 GitHub 上找到：github.com/PacktPublishing/Learn-CUDA-Programming。

本章的代码示例是使用 CUDA Toolkit 的 10.1 版本开发和测试的，但建议尽可能使用最新的 CUDA 版本。

来自 CUDA 的 Hello World

CUDA 是一个包括 CPU 和 GPU 在内的异构编程模型。CUDA C/C++编程接口由 C 语言扩展组成，以便您可以将源代码的部分目标定为在设备（GPU）上并行执行。它基于行业标准的 C/C++，并提供了一系列 C 函数库，可以在主机（CPU）上执行，以便它可以与设备进行交互。

在 CUDA 中，有两个相互配合的处理器。主机通常被称为 CPU，而设备通常被称为 GPU。主机负责调用设备函数。正如我们已经提到的，运行在 GPU 上的代码的一部分被称为设备代码，而在 CPU 上运行的串行代码被称为主机代码。

让我们从在 C 中编写我们的第一个 CUDA 代码开始。我们的意图是采取一个系统化的逐步方法，从一些顺序代码开始，通过添加一些额外的关键字将其转换为 CUDA 感知代码。正如我们之前提到的，没有必要学习一门新语言，我们只需要在现有语言中添加一些关键字，以便在 CPU 和 GPU 的异构环境中运行它。

让我们来看看我们的第一段代码。这段代码的作用只是从主机和设备上打印 Hello World!

#include<stdio.h>
#include<stdlib.h>

__global__ void print_from_gpu(void) {
    printf("Hello World! from thread [%d,%d] \
        From device\n", threadIdx.x,blockIdx.x);
}

int main(void) {
    printf("Hello World from host!\n");
    print_from_gpu<<<1,1>>>();
    cudaDeviceSynchronize();
    return 0;
}

让我们尝试编译和运行前面的代码片段：

编译代码：将前面的代码放入一个名为hello_world.cu的文件中，并使用NVIDIA C Compiler（nvcc）进行编译。请注意，文件的扩展名是.cu，这告诉编译器这个文件里面有 GPU 代码：

$ nvcc -o hello_world hello_world.cu

执行 GPU 代码：在执行 GPU 代码后，我们应该收到以下输出：

到目前为止，您可能已经注意到 CUDA C 代码的使用方式并没有太大不同，只需要学习一些额外的构造来告诉编译器哪个函数是 GPU 代码，以及如何调用 GPU 函数。这并不像我们需要完全学习一门新语言。

在前面的代码中，我们添加了一些构造和关键字，如下：

__global__：在函数之前添加此关键字，告诉编译器这是一个将在设备上而不是主机上运行的函数。但请注意，它是由主机调用的。这里另一个重要的事情是设备函数的返回类型始终是"void"。算法的数据并行部分在设备上作为内核执行。
<<<,>>>: 这个关键字告诉编译器这是对设备函数的调用，而不是对主机函数的调用。此外，1,1参数基本上决定了在内核中启动的线程数。我们将稍后介绍尖括号内的参数。目前，1,1参数基本上意味着我们只启动一个线程的内核，也就是说，除了打印之外，我们在代码中没有做任何重要的事情。
threadIdx.x, blockIdx.x: 这是给所有线程的唯一 ID。我们将在下一节更详细地介绍这个主题。
cudaDeviceSynchronize(): CUDA 中的所有内核调用都是异步的。在调用内核后，主机变得空闲，并在之后开始执行下一条指令。这应该不足为奇，因为这是一个异构环境，因此主机和设备都可以并行运行，以利用可用的处理器类型。如果主机需要等待设备完成，CUDA 编程提供了 API 使主机代码等待设备函数完成。其中一个 API 是cudaDeviceSynchronize，它会等待所有先前对设备的调用完成。

尝试删除cudaDeviceSynchronize()调用，看看设备输出是否可见。或者，尝试在打印主机代码之前放置这个调用。

线程层次结构

现在，让我们开始玩弄两个参数，即threadIdx.x和blockIdx.x。

实验 1：首先，将参数从<<<1,1>>>更改为<<<2,1>>>并查看输出。运行多个线程-单个块的 Hello World 代码的输出应该如下：

正如我们所看到的，现在我们不是一个线程，而是两个线程打印值。请注意，它们的唯一 ID 是不同的。

实验 2：现在，不要更改第一个参数，而是更改第二个参数，即将<<<1,1>>>更改为<<<1,2>>>，并观察运行多个单线程块的 Hello World 代码的输出，如下所示：

如您所见，被启动到内核中的线程总数是两个，就像以前一样——唯一的区别是它们的 ID 不同。那么，这些线程和块的概念是什么？为了解决这个问题，让我们更深入地了解 GPU 架构。

GPU 架构

CUDA 变得如此受欢迎的一个关键原因是因为硬件和软件被设计和紧密绑定，以获得应用程序的最佳性能。因此，有必要展示软件 CUDA 编程概念与硬件设计本身之间的关系。

以下截图显示了 CUDA 的两个方面：

我们可以看到，CUDA 软件已经映射到了 GPU 硬件。

根据前面的截图，以下表解释了 CUDA 编程模型的软件和硬件映射：

软件	执行于/作为	硬件
CUDA 线程	CUDA 核心/SIMD 代码
CUDA 块	流多处理器
网格/内核	GPU 设备

让我们详细看一下前表的组件：

CUDA 线程：CUDA 线程在 CUDA 核心上执行。CUDA 线程不同于 CPU 线程。CUDA 线程非常轻量级，并提供快速的上下文切换。快速上下文切换的原因是由于 GPU 中有大量的寄存器和硬件调度器。线程上下文存在于寄存器中，而不是像 CPU 中那样在较低的内存层次结构中，比如缓存中。因此，当一个线程处于空闲/等待状态时，另一个准备好的线程几乎可以立即开始执行。每个 CUDA 线程必须执行相同的内核，并且独立地处理不同的数据（SIMT）。
CUDA 块：CUDA 线程被组合成一个称为 CUDA 块的逻辑实体。CUDA 块在单个流多处理器（SM）上执行。一个块在一个 SM 上运行，也就是说，一个块内的所有线程只能在一个 SM 的核心上执行，不会在其他 SM 的核心上执行。每个 GPU 可能有一个或多个 SM，因此为了有效地利用整个 GPU，用户需要将并行计算划分为块和线程。
GRID/核心：CUDA 块被组合成一个称为 CUDA GRID 的逻辑实体。然后在设备上执行 CUDA GRID。

乍一看，这可能听起来有些复杂。在接下来的部分，我们将以向量加法的示例来解释这个问题。希望事情会变得更清晰。

使用 CUDA 进行向量加法

我们要解决的问题是向量加法。正如我们所知，向量加法是一种数据并行操作。我们的数据集包括三个数组：A、B和C。每个元素执行相同的操作：

Cx = Ax + Bx

每个加法是相互独立的，但所有 CUDA 线程都执行相同的操作。要开始，根据以下步骤配置您的环境：

准备你的 GPU 应用程序。这段代码将放在01_cuda_introduction/01_vector_addition中。
使用nvcc编译器编译您的应用程序，命令如下：

$nvcc -o vector_addition vector_addition.cu

上述代码是顺序代码。我们将按照以下步骤逐步将此代码转换为可以在 GPU 上运行的代码：

#include<stdio.h>
#include<stdlib.h>

#define N 512

void host_add(int *a, int *b, int *c) {
    for(int idx=0;idx<N;idx++)
        c[idx] = a[idx] + b[idx];
}

//basically just fills the array with index.
void fill_array(int *data) {
    for(int idx=0;idx<N;idx++)
        data[idx] = idx;
}

void print_output(int *a, int *b, int*c) {
    for(int idx=0;idx<N;idx++)
        printf("\n %d + %d = %d", a[idx] , b[idx], c[idx]);
}

int main(void) {
    int *a, *b, *c;
    int size = N * sizeof(int);
   // Alloc space for host copies of a, b, c and setup input values
    a = (int *)malloc(size); fill_array(a);
    b = (int *)malloc(size); fill_array(b);
    c = (int *)malloc(size);
    host_add(a,b,c);
    print_output(a,b,c);
    free(a); free(b); free(c);
    return 0;
}

在转换顺序代码之前，让我们看一下 CUDA 代码和顺序代码之间所采取的基本变化或步骤：

顺序代码	CUDA 代码
步骤 1	在 CPU 上分配内存，即`malloc new`。
步骤 2	填充/初始化 CPU 数据。
步骤 3	调用处理数据的 CPU 函数。在这种情况下，实际算法是向量加法。
步骤 4	处理数据，这里是打印出来的。
步骤 5	用`<<<,>>>`括号调用 GPU 函数。
步骤 6	用`cudaDeviceSynchronize`同步设备和主机。
步骤 7	用`cudaMemcpy`将数据从设备传输到主机。
步骤 8	处理数据，这里是打印出来的。

本书不是 CUDA API 指南的替代品，也不涵盖所有 CUDA API。如需广泛使用 API，请参考 CUDA API 指南。

正如我们所看到的，CUDA 处理流程有一些额外的步骤需要添加到顺序代码中。具体如下：

在 GPU 上分配内存：CPU 内存和 GPU 内存是物理上分开的内存。malloc在 CPU 的 RAM 上分配内存。GPU 核心/设备函数只能访问已分配/指向设备内存的内存。要在 GPU 上分配内存，我们需要使用cudaMalloc API。与malloc命令不同，cudaMalloc不会返回指向已分配内存的指针；相反，它以指针引用作为参数，并更新相同的已分配内存。
将数据从主机内存传输到设备内存：主机数据然后被复制到使用前一步中使用的cudaMalloc命令分配的设备内存。用于在主机和设备之间复制数据的 API 是cudaMemcpy。与其他memcopy命令一样，此 API 需要目标指针、源指针和大小。它额外需要一个参数，即复制的方向，也就是说，我们是从主机到设备复制，还是从设备到主机复制。在 CUDA 的最新版本中，这是可选的，因为驱动程序能够理解指针是指向主机内存还是设备内存。请注意，cudaMemcpy有一个异步的替代方案。这将在其他章节中更详细地介绍。
调用和执行 CUDA 函数：如同在 Hello World CUDA 程序中所示，我们通过<<<,>>>括号调用一个核函数，它提供了块和线程大小的参数。在完成所有步骤后，我们将更详细地介绍这一点。
同步：正如我们在 Hello World 程序中提到的，核函数调用是异步的。为了确保主机确保核执行已完成，主机调用cudaDeviceSynchronize函数。这确保了之前启动的所有设备调用都已完成。
将数据从主机内存传输到设备内存：使用相同的cudaMemcpy API 将数据从设备复制回主机，用于后处理或验证任务，如打印。与第一步相比，唯一的变化是我们颠倒了复制的方向，也就是说，目标指针指向主机，而源指针指向在内存中分配的设备。
释放分配的 GPU 内存：最后，使用cudaFree API 释放分配的 GPU 内存。

更改顺序向量加法代码的main函数以反映这些新步骤。main函数将如下所示：

int main(void) {
    int *a, *b, *c;
    int *d_a, *d_b, *d_c; // device copies of a, b, c
    int size = N * sizeof(int);

    // Alloc space for host copies of a, b, c and setup input values
    a = (int *)malloc(size); fill_array(a);
    b = (int *)malloc(size); fill_array(b);
    c = (int *)malloc(size);

    // Alloc space for device copies of vector (a, b, c)
    cudaMalloc((void **)&d_a, N* * sizeof(int));
    cudaMalloc((void **)&d_b, N* *sizeof(int));
    cudaMalloc((void **)&d_c, N* * sizeof(int));

    // Copy from host to device
    cudaMemcpy(d_a, a, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, N* sizeof(int), cudaMemcpyHostToDevice);

    device_add<<<1,1>>>(d_a,d_b,d_c);

    // Copy result back to host
    cudaMemcpy(c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);

    print_output(a,b,c);
    free(a); free(b); free(c);

    //free gpu memory
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);

    return 0;
}

现在，让我们看一下如何编写核心代码并管理线程和块大小。为此，我们将进行多个实验。

实验 1 - 创建多个块

在这一部分，我们将利用 CUDA 块在 GPU 上并行运行向量加法代码。将会暴露与我们如何索引 CUDA 块相关的附加关键字。更改对device_add函数的调用如下：

//changing from device_add<<<1,1>>> to
device_add<<<N,1>>>

这将使device_add函数并行执行N次，而不是一次。device_add函数的每个并行调用称为一个块。现在，让我们添加一个__global__设备函数，如下所示：

__global__ void device_add(int *a, int *b, int *c) {
 c[blockIdx.x] = a[blockIdx.x] + b[blockIdx.x];
}

通过使用blockIdx.x来索引数组，每个块处理数组的不同元素。在设备上，每个块可以并行执行。让我们看一下以下的屏幕截图：

前面的屏幕截图代表了向量加法 GPU 代码，其中每个块显示了多个单线程块的索引。

实验 2 - 创建多个线程

在这一部分，我们将利用 CUDA 线程在 GPU 上并行运行向量加法代码。将会暴露与我们如何索引 CUDA 线程相关的附加关键字。

一个块可以分成多个线程。更改对device_add函数的调用如下：

//changing from device_add<<<1,1>>> to
device_add<<<1,N>>>

这将并行执行device_add函数N次，而不是一次。device_add函数的每个并行调用被称为一个线程。更改设备例程以反映内核，如下所示：

__global__ void device_add(int *a, int *b, int *c) {
     c[threadIdx.x] = a[threadIdx.x] + b[threadIdx.x];
}

一个显著的区别是，我们使用threadIdx.x而不是blockIdx.x，如下面的截图所示：

前面的截图代表了向量加法 GPU 代码，其中每个块显示了单个块-多个线程的索引。

实验 3 - 结合块和线程

到目前为止，我们已经通过在实验 1 - 创建多个块部分使用多个块和一个线程，以及在实验 2 - 创建多个线程部分使用一个块和多个线程来查看并行向量加法。在这个实验中，我们将使用多个块以及包含多个线程的单独块。在如何找到索引方面，这变得更具挑战性，因为我们需要结合threadIdx和blockIdx来生成一个唯一的 ID。

让我们看一下两种不同组合的场景，开发人员可以从中选择：

场景 1： 假设向量元素的总数是 32。每个块包含八个线程，总共有四个块。
场景 2： 假设向量元素的总数是 32。每个块包含四个线程，总共有八个块。

在这两种情况下，并行执行的数量都是 32，所有 32 个元素都会并行填充。开发人员根据问题的大小和每个硬件的限制选择块内的线程和块的数量。我们将在另一章节中详细介绍基于架构的正确尺寸选择的细节。

以下截图显示了不同块和线程配置的向量加法 GPU 索引代码：

现在，让我们看看如何改变内核代码以结合线程和块来计算全局索引：

__global__ void device_add(int *a, int *b, int *c) {
     int index = threadIdx.x + blockIdx.x * blockDim.x;
     c[index] = a[index] + b[index];
}

在从main()函数中调用内核时，开发人员选择了块和线程的配置，如前面提到的两种情况所示的代码：

场景 1： 以下是用于计算每个块八个线程的向量加法 GPU 网格和块大小的代码：

threads_per_block = 8;
no_of_blocks = N/threads_per_block;
device_add<<<no_of_blocks,threads_per_block>>>(d_a,d_b,d_c);

场景 2： 以下是用于计算每个块四个线程的向量加法 GPU 网格和块大小的代码：

threads_per_block = 4;
no_of_blocks = N/threads_per_block;
device_add<<<no_of_blocks,threads_per_block>>>(d_a,d_b,d_c);

通过线程和块的组合，可以计算出线程的唯一 ID。如前面的代码所示，所有线程都被赋予另一个变量。这被称为blockDim。这个变量包含了块的维度，也就是每个块的线程数。让我们看一下下面的截图：

在这里，我们可以看到场景 1 的向量加法 GPU 索引计算。

为什么要费心处理线程和块？

可能不明显为什么我们需要这种额外的线程和块的层次结构。它增加了开发人员需要找到正确的块和网格大小的复杂性。全局索引也变得具有挑战性。这是因为 CUDA 编程模型设置了限制。

与并行块不同，线程有有效的通信和同步机制。现实世界的应用程序需要线程之间进行通信，并且可能希望在继续之前等待某些数据进行交换。这种操作需要线程进行通信，CUDA 编程模型允许同一块内的线程进行通信。属于不同块的线程在内核执行期间无法进行通信/同步。这种限制允许调度程序独立地在 SM 上调度块。其结果是，如果发布了具有更多 SM 的新硬件，并且代码具有足够的并行性，则代码可以线性扩展。换句话说，这允许硬件根据 GPU 的能力并行运行块的数量。

线程之间使用一种称为共享内存的特殊内存进行通信。我们将在第二章中广泛介绍共享内存，即CUDA 内存管理，在那里我们将介绍 GPU 中的其他内存层次结构及其最佳使用方法。以下屏幕截图演示了在不同 GPU 上扩展块，这些 GPU 包含不同数量的 SM：

现在，让我们更多地了解在多个维度中启动内核。

在多个维度中启动内核

到目前为止，我们一直在一维中启动线程和块。这意味着我们只使用一个维度的索引；例如，我们一直在使用threadIdx.x，其中x表示我们只使用一个x维度的线程索引。同样，我们一直在使用blockIdx.x，其中x表示我们只使用一个x维度的块索引。我们可以在一、二或三个维度中启动线程和块。在二维中启动线程和块的一个例子是当我们在图像上进行并行操作，例如使用滤波器模糊图像。开发人员可以选择在二维中启动线程和块，这是一个更自然的选择，因为图像在本质上是二维的。

重要的是要了解每个 GPU 架构也对线程和块的维度施加了限制。例如，NVIDIA Pascal 卡允许在x和y维度中每个线程块最多有 1,024 个线程，而在z维度中，您只能启动 64 个线程。同样，在 Pascal 架构中，网格中的最大块数限制为y和z维度中的 65,535 个，x维度中为2³¹ -1。如果开发人员使用不受支持的维度启动内核，应用程序会抛出运行时错误。

到目前为止，我们一直假设我们编写的代码是没有错误的。但在现实世界中，每个程序员都会写有错误的代码，必须捕捉这些错误。在下一节中，我们将看看 CUDA 中的错误报告是如何工作的。

CUDA 中的错误报告

在 CUDA 中，主机代码管理错误。大多数 CUDA 函数调用cudaError_t，它基本上是一个枚举类型。cudaSuccess（值 0）表示0错误。用户还可以使用cudaGetErrorString()函数，该函数返回描述错误条件的字符串。

 cudaError_t e;
 e = cudaMemcpy(...);
 if(e)
     printf("Error: %sn", cudaGetErrorString(err));

内核启动没有返回值。我们可以在这里使用cudaGetLastError()这样的函数，它返回最后一个 CUDA 函数（包括内核启动）的错误代码。在多个错误的情况下，只报告最后一个错误：

MyKernel<<< ... >>> (...);
cudaDeviceSynchronize();
e = cudaGetLastError();

在生产代码中，建议在逻辑检查点处使用错误检查代码，因为即使 GPU 内核崩溃，CPU 代码也会继续正常执行，导致结果不正确。

在下一节中，我们将向您介绍 CUDA 编程模型中支持的数据类型。

CUDA 中的数据类型支持

与任何处理器架构一样，GPU 也有不同类型的内存，每种用于不同的目的。我们将在第二章中更详细地介绍它们，CUDA 内存管理。然而，重要的是要理解支持的不同数据类型及其对性能和精度的影响。CUDA 编程支持开发人员在其各自语言中熟悉的所有标准数据类型。除了不同大小的标准数据类型（char为 1 字节，float为 4 字节，double为 8 字节等）之外，它还支持矢量类型，如float2和float4。

建议数据类型自然对齐，因为对于大小为 1、2、4、8 或 16 字节的数据类型的对齐数据访问，可以确保 GPU 调用单个内存指令。如果它们没有对齐，编译器将生成多个交错的指令，导致内存和指令总线的低效利用。因此，建议在 GPU 内存中使用自然对齐的类型。对于char、short、int、long、long long、float和double等内置类型，如float2和float4，对齐要求会自动满足。

此外，CUDA 编程还支持复杂的数据结构，如结构和类（在 C 和 C++的上下文中）。对于复杂的数据结构，开发人员可以利用对齐说明符来强制编译器满足对齐要求，如下面的代码所示：

struct __align__(16) {
    float r;
    float g;
    float b;
};

每个 GPU 都有一组有限的核心，因此 FLOPS 是不同的。例如，具有 Volta 架构的 Tesla V100 卡具有 2560 个 FP64 核心（双精度），而具有两倍数量的 32 位单精度核心。很明显，根据算法的精度要求使用正确的数据类型是至关重要的。现在正在开发混合精度算法，以利用不同类型的核心，其中算法的某些部分以更高精度运行，而某些部分以较低精度运行。我们将在即将到来的章节中更多地涵盖这个主题。目前，重要的是要理解 GPU 内存层次结构是不同的，因此使用正确的数据类型很重要。

虽然这是对 GPU 支持的数据类型的一般介绍，但有关所有支持的数据类型的更多详细信息可以在docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#built-in-vector-types找到。

总结

在本章中，我们通过历史和高性能计算为您提供了异构计算的视角。我们详细介绍了两个处理器，即 CPU 和 GPU 的不同之处。我们还在 GPU 上编写了一个 Hello World 和矢量加法 CUDA 程序。最后，我们看了如何检测 CUDA 中的错误，因为对 CUDA API 的大多数调用都是异步的。

在下一章中，我们将看一下不同类型的 GPU 内存以及如何最优地利用它们。

第二章：CUDA 内存管理

正如我们在第一章中所描述的，CUDA 编程简介，CPU 和 GPU 架构在根本上是不同的，它们的内存层次结构也是不同的。它们不仅在大小和类型上有所不同，而且在目的和设计上也有所不同。到目前为止，我们已经学习了每个线程如何通过索引（blockIdx和threadIdx）访问自己的数据。我们还使用了诸如cudaMalloc之类的 API 在设备上分配内存。GPU 中有许多内存路径，每个路径的性能特征都不同。启动 CUDA 核心可以帮助我们实现最大性能，但只有在以最佳方式使用正确类型的内存层次结构时才能实现。将数据集映射到正确的内存类型是开发人员的责任。

根据经验，如果我们绘制一个图表，概述 GPU 上的顶级应用性能约束，它将看起来像以下图表：

上述饼图粗略地分解了大多数基于 CUDA 的应用程序中出现的性能问题。很明显，大多数情况下，应用程序的性能将受到与内存相关的约束的限制。根据应用程序和采取的内存路径，内存相关的约束进一步划分。

让我们以不同的方式来看待这种方法，并了解有效使用正确类型的内存的重要性。最新的 NVIDIA GPU 采用 Volta 架构，提供了 7,000 GFLOP 的峰值性能，其设备内存带宽为 900 GB/s。您将首先注意到的是 FLOP 与内存带宽的比率，大约为 7:1。这是假设所有线程都访问 4 字节（浮点数）数据执行操作。执行此操作所需的总带宽是47,000 = 28,000* GB/s，即达到峰值性能所需的带宽。900 GB/s 将执行限制为 225 GFLOP。这将执行速率限制为峰值的 3.2%（225 GFLOP 是设备的 7,000 GFLOP 峰值的 3.2%）。正如您现在所知，GPU 是一种隐藏延迟的架构，有许多可用于执行的线程，这意味着它在理论上可以容忍长的内存访问延迟。然而，对内存的过多调用可能会导致一些 SMs 空闲，导致一些线程停顿或等待。CUDA 架构提供了其他几种方法，我们可以使用这些方法来访问内存，以解决内存瓶颈问题。

从 CPU 内存到被 SM 用于处理的数据路径在下图中展示。在这里，我们可以看到数据元素在到达 SM 核心进行计算之前的旅程。每个内存带宽的数量级都不同，访问它们的延迟也不同：

在上图中，我们可以看到从 CPU 到达寄存器的数据路径，最终计算是由 ALU/核心完成的。

下图显示了最新 GPU 架构中存在的不同类型的内存层次结构。每种内存可能具有不同的大小、延迟、吞吐量和应用程序开发人员的可见性：

上图显示了最新 GPU 架构中存在的不同类型的内存及其在硬件中的位置。

在本章中，您将学习如何最佳地利用不同类型的 GPU 内存。我们还将研究 GPU 统一内存的最新特性，这使得程序员的生活变得更简单。本章将详细介绍以下内存主题：

全局内存/设备内存
共享内存
只读数据/缓存
固定内存
统一内存

但在我们查看内存层次结构之前，我们将遵循优化周期，如下所示：

步骤 1：分析
步骤 2：并行化
步骤 3：优化

对应用程序的分析要求我们不仅要了解我们应用程序的特性，还要了解它在 GPU 上的有效运行方式。为此，我们将首先向您介绍 Visual Profiler，然后再进入 GPU 内存。由于我们在这里使用了一些最新的 CUDA 功能，请在继续本章之前阅读以下部分。

技术要求

本章需要一台带有现代 NVIDIA GPU（Pascal 架构或更高版本）的 Linux PC，以及安装了所有必要的 GPU 驱动程序和 CUDA Toolkit（10.0 或更高版本）。如果您不确定您的 GPU 架构，请访问 NVIDIA GPU 网站developer.nvidia.com/cuda-gpus进行确认。本章的代码也可以在 GitHub 上找到github.com/PacktPublishing/Learn-CUDA-Programming。

本章的示例代码示例是使用 CUDA Toolkit 的 10.1 版本开发和测试的。但是，建议使用最新的 CUDA 版本或更高版本。

在下一节中，我们将向您介绍 Visual Profiler，它将帮助我们分析我们的应用程序。我们还将看一下它在 GPU 上的运行情况。

NVIDIA Visual Profiler

为了了解不同内存层次结构的有效利用，重要的是在运行时分析应用程序的特性。分析器是非常方便的工具，可以测量和显示不同的指标，帮助我们分析内存、SM、核心和其他资源的使用方式。 NVIDIA 决定提供一个 API，供分析器工具的开发人员用于连接到 CUDA 应用程序，随着时间的推移，一些分析工具已经发展出来，如 TAU 性能系统、Vampir Trace 和 HPC Toolkit。所有这些工具都利用CUDA 分析器工具接口（CUPTI）为 CUDA 应用程序提供分析信息。

NVIDIA 本身开发并维护作为 CUDA Toolkit 的一部分提供的分析工具。本章使用这两个分析工具（NVPROF 和 NVVP）来演示不同内存类型的有效使用，并不是分析工具的指南。

我们将使用 NVPROF 或 NVVP 来演示 CUDA 应用程序的特性。NVPROF 是一个命令行工具，而nvvp具有可视化界面。nvvp有两种格式，一种是独立版本，另一种是集成在 Nsisght Eclipse 中的版本。

我们将广泛使用的 NVVP 分析器窗口如下所示：

这是在 macOS 上拍摄的 NVVP 9.0 版本窗口快照。

窗口中有四个视图可用：时间轴、指南、分析结果和摘要。时间轴视图显示了随时间发生的 CPU 和 GPU 活动。Visual Profiler 显示了 CUDA 编程模型的内存层次结构的摘要视图。分析视图显示了分析结果。Visual Profiler 提供了两种分析模式：

引导分析：顾名思义，它通过逐步方法指导开发人员了解关键性能限制器。我们建议初学者在成为了解不同指标的专家之前先使用此模式，然后再转到无引导模式。
无引导分析：开发人员必须手动查看此模式下的结果，以了解性能限制器。

CUDA Toolkit 提供了两个 GPU 应用程序性能分析工具，NVIDIA Profiler（NVPROF）和NVIDIA Visual Profiler（NVVP）。为了获得性能限制器信息，我们需要进行两种类型的分析：时间线分析和度量分析。此代码可在 02_memory_overview/04_sgemm 中访问。分析命令可以执行如下：

$ nvcc -o sgemm sgemm.cu
$ nvprof -o sgemm.nvvp ./sgemm
$ nvprof --analysis-metrics -o sgemm-analysis.nvvp ./sgemm

让我们打开 Visual Profiler。如果你使用的是 Linux 或 OSX，你可以在终端中执行 nvvp。或者，你可以从安装了 CUDA Toolkit 的二进制文件中找到 nvvp 可执行文件。如果你使用的是 Windows，你可以使用 Windows 搜索框执行该工具，命令为 nvvp。

要打开两个分析数据，我们将使用“文件”|“导入...”菜单，如下所示：

然后，我们将继续点击底部的“下一步”按钮：

。

我们的 CUDA 应用程序使用一个进程。因此，让我们继续点击底部的“下一步”按钮：

现在，让我们将收集的分析数据放入 Visual Profiler 中。以下截图显示了一个示例。通过右侧的“浏览...”按钮，将时间线数据放入第二个文本框。然后，以相同的方式将度量分析数据放入下一个文本框中：

有关性能分析工具的详细使用，请参阅 CUDA Profiling 指南，该指南作为 CUDA Toolkit 的一部分提供（相应的网页链接为 docs.nvidia.com/cuda/profiler-users-guide/index.html）。

在基于 Windows 的系统中，在安装了 CUDA Toolkit 后，你可以从“开始”菜单中启动 Visual Profiler。在具有 X11 转发的 Linux 系统中，你可以通过运行 nvvp 命令来启动 Visual Profiler，nvvp 代表 NVIDIA Visual Profiler：

$ ./nvvp

既然我们现在对将要使用的分析工具有了一个公平的理解，让我们进入第一个也是绝对最关键的 GPU 内存——全局内存/设备内存。

全局内存/设备内存

本节将详细介绍如何使用全局内存，也称为设备内存。在本节中，我们还将讨论如何高效地将数据从全局内存加载/存储到缓存中。由于全局内存是一个暂存区，所有数据都从 CPU 内存中复制到这里，因此必须充分利用这种内存。全局内存或设备内存对于内核中的所有线程都是可见的。这种内存也对 CPU 可见。

程序员使用 cudaMalloc 和 cudaFree 显式地管理分配和释放。数据使用 cudaMalloc 分配，并声明为 __device__。全局内存是从 CPU 使用 cudaMemcpy API 传输的所有内存的默认暂存区。

全局内存上的矢量加法

我们在第一章中使用的矢量加法示例演示了全局内存的使用。让我们再次查看代码片段，并尝试理解全局内存的使用方式：

__global__ void device_add(int *a, int *b, int *c) {
     int index = threadIdx.x + blockIdx.x * blockDim.x;
     c[index] = a[index] + b[index];
}
int main (void) {
...
    // Alloc space for device copies of a, b, c
    cudaMalloc((void **)&d_a, size);
    cudaMalloc((void **)&d_b, size);
    cudaMalloc((void **)&d_c, size);
...

   // Free space allocated for device copies
   cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
...

}

cudaMalloc 在设备内存上分配数据。内核中的参数指针（a、b 和 c）指向这个设备内存。我们使用 cudaFree API 释放这个内存。正如你所看到的，块中的所有线程都可以在内核中访问这个内存。

此代码可在 02_memory_overview/01_vector_addition 中访问。要编译此代码，可以使用以下命令：

$ nvcc -o vec_addition ./vector_addition_gpu_thread_block.cu

这是一个使用全局内存的简单示例。在下一节中，我们将看看如何最优地访问数据。

合并与未合并的全局内存访问

为了有效使用全局内存，了解 CUDA 编程模型中 warp 的概念是非常重要的，这是我们到目前为止忽略的。warp 是 SM 中的线程调度/执行单位。一旦一个块被分配给一个 SM，它被划分为一个 32 个线程的单位，称为warp。这是 CUDA 编程中的基本执行单位。

为了演示 warp 的概念，让我们看一个例子。如果两个块被分配给一个 SM，每个块有 128 个线程，那么块内的 warp 数量是128/32 = 4个 warp，SM 上的总 warp 数量是4 * 2 = 8个 warp。以下图表显示了 CUDA 块如何在 GPU SM 上被划分和调度：

块和 warp 在 SM 和其核心上的调度更多地是与体系结构相关的，对于 Kepler、Pascal 和最新的 Volta 等不同的架构，情况会有所不同。目前，我们可以忽略调度的完整性。在所有可用的 warp 中，具有下一条指令所需操作数的 warp 变得可以执行。根据运行 CUDA 程序的 GPU 的调度策略，选择要执行的 warp。当被选择时，warp 中的所有线程执行相同的指令。CUDA 遵循单指令，多线程（SIMT）模型，也就是说，warp 中的所有线程在同一时间实例中获取和执行相同的指令。为了最大程度地利用全局内存的访问，访问应该合并。合并和未合并之间的区别如下：

合并的全局内存访问： 顺序内存访问是相邻的。
未合并的全局内存访问： 顺序内存访问不是相邻的。

以下图表更详细地展示了这种访问模式的示例。图表的左侧显示了合并访问，其中 warp 中的线程访问相邻数据，因此导致了一个 32 位宽的操作和 1 次缓存未命中。图表的右侧显示了一种情况，即 warp 内的线程访问是随机的，可能导致调用 32 次单个宽度的操作，因此可能有 32 次缓存未命中，这是最坏的情况：

为了进一步理解这个概念，我们需要了解数据如何通过缓存行从全局内存到达。

情况 1： warp 请求 32 个对齐的、4 个连续的字节

地址落在 1 个缓存行内和一个 32 位宽的操作内。总线利用率为 100%，也就是说，我们利用从全局内存中获取的所有数据到缓存中，并没有浪费任何带宽。如下图所示：

上图显示了合并的访问，导致了总线的最佳利用。

情况 2： warp 请求 32 个分散的 4 字节单词

虽然 warp 需要 128 字节，但在未命中时执行了 32 次单个宽度的获取，导致32 * 128字节在总线上移动。如下图所示，总线利用率实际上低于 1%：

上图显示了未合并的访问，导致了总线带宽的浪费。

正如我们在前面的图表中看到的，warp 内的线程如何从全局内存中访问数据非常重要。为了最大程度地利用全局内存，改善合并是非常重要的。有多种可以使用的策略。其中一种策略是改变数据布局以改善局部性。让我们看一个例子。将滤波器应用于图像或将掩模应用于图像的计算机视觉算法需要将图像存储到数据结构中。当开发人员声明图像类型时，有两种选择。

以下代码片段使用Coefficients_SOA数据结构以数组格式存储数据。Coefficients_SOA结构存储与图像相关的数据，如 RGB、色调和饱和度值：

//Data structure representing an image stored in Structure of Array Format
struct Coefficients_SOA {
 int r;
 int b;
 int g;
 int hue;
 int saturation;
 int maxVal;
 int minVal;
 int finalVal;
};

以下图表显示了关于Coefficients_SOA存储数据的数据布局，以及在内核中由不同线程访问数据的情况：

通过这样做，我们可以看到 AOS 数据结构的使用导致了不连续的全局内存访问。

同样的图像可以以数组结构格式存储，如下面的代码片段所示：

//Data structure representing an image stored in Array of Structure Format
struct Coefficients_AOS {
 int* r;
 int* b;
 int* g;
 int* hue;
 int* saturation;
 int* maxVal;
 int* minVal;
 int* finalVal;
};

以下图表显示了关于Coefficients_AOS存储数据的数据布局，以及在内核中由不同线程访问数据的情况：

通过这样做，我们可以看到使用 SOA 数据结构导致了不连续的全局内存访问。

虽然 CPU 上的顺序代码更喜欢 AOS 以提高缓存效率，但在单指令多线程（SIMT）模型（如 CUDA）中，SOA 更受欢迎，以提高执行和内存效率。

让我们尝试通过使用分析器来分析这一方面。根据以下步骤配置你的环境：

准备好你的 GPU 应用程序。例如，我们将使用两段代码来演示全局内存的有效使用。aos_soa.cu文件包含了使用 AOS 数据结构的朴素实现，而aos_soa_solved.cu则使用了 SOA 数据结构，可以有效地利用全局内存。这段代码可以在02_memory_overview/02_aos_soa中找到。
使用nvcc编译器编译你的应用程序，然后使用nvprof编译器对其进行分析。以下命令是对此的一个nvcc命令的示例。然后我们使用nvprof命令对应用程序进行分析。还传递了--analysis-metrics标志，以便我们可以获得内核的指标。
生成的分析文件，即aos_soa.prof和aos_soa_solved.prof，然后加载到 NVIDIA Visual Profiler 中。用户需要从“文件|打开”菜单中加载分析输出。此外，不要忘记在文件名选项中选择“所有文件”。

$ nvcc -o aos_soa ./aos_soa.cu
$ nvcc -o aos_soa_solved ./aos_soa_solved.cu
$ nvprof --analysis-metrics --export-profile aos_soa.prof ./aos_soa
$ nvprof --analysis-metrics --export-profile aos_soa_solved.prof ./aos_soa_solved

以下是分析输出的屏幕截图。这是一个使用 AOS 数据结构的朴素实现：

前面的图表显示了在引导分析模式下分析器的输出。

你将看到的第一件事是分析器明确指出应用程序受到内存限制。正如你所看到的，分析器不仅显示指标，还分析了这些指标的含义。在这个例子中，由于我们使用 AOS，分析器明确指出访问模式不高效。但编译器是如何得出这个结论的呢？让我们看一下下面的屏幕截图，它提供了更多的细节：

正如我们所看到的，它清楚地说明了访问数据的理想事务数为四，而实际运行时进行了 32 次事务/访问。

当我们将数据结构从 AOS 更改为 SOA 时，瓶颈得到了解决。当你运行aos_soa_solved可执行文件时，你会发现内核时间减少了，这对我们的计时来说是一个改进。在 V100 16 GB 卡上，时间从 104 微秒减少到 47 微秒，这是一个2.2x的加速因子。分析输出aos_soa_solved.prof将显示内核仍然受到内存限制，这是非常明显的，因为我们读写的内存数据比进行计算时要多。

内存吞吐量分析

对于应用程序开发人员来说，了解应用程序的内存吞吐量非常重要。这可以通过两种方式来定义：

从应用程序的角度来看： 计算应用程序请求的字节数
从硬件的角度来看： 计算硬件传输的字节数

这两个数字完全不同。其中许多原因包括未协调的访问导致未利用所有事务字节，共享内存银行冲突等。我们应该从内存角度使用两个方面来分析应用程序：

地址模式：在实际代码中确定访问模式是非常困难的，因此使用诸如性能分析器之类的工具变得非常重要。性能分析器显示的指标，如全局内存效率和每次访问的 L1/L2 事务，需要仔细观察。
飞行中的并发访问数量：由于 GPU 是一种隐藏延迟的架构，饱和内存带宽变得非常重要。但是确定并发访问的数量通常是不够的。此外，从硬件的角度来看，吞吐量与理论值相比要不同得多。

下图演示了每个 SM 中飞行的约 6KB 数据可以达到 Volta 架构峰值带宽的 90%。在以前的一代架构上进行相同的实验会得到不同的图表。一般来说，建议了解特定架构的 GPU 内存特性，以便从硬件中获得最佳性能：

本节为我们提供了全局内存的示例用法以及如何以最佳方式利用它。有时，全局内存的协调数据访问很困难（例如，在 CFD 领域，对于非结构化网格，相邻单元格的数据可能不会相邻存储在内存中）。为了解决这样的问题或减少对性能的影响，我们需要利用另一种形式的内存，称为共享内存。

共享内存

共享内存一直在 CUDA 内存层次结构中扮演着重要角色，被称为用户管理的缓存。这为用户提供了一种机制，可以以协调的方式从全局内存中读取/写入数据并将其存储在内存中，这类似于缓存但可以由用户控制。在本节中，我们不仅将介绍利用共享内存的步骤，还将讨论如何有效地从共享内存中加载/存储数据以及它在银行中的内部排列。共享内存只对同一块中的线程可见。块中的所有线程看到共享变量的相同版本。

共享内存具有类似于 CPU 缓存的好处；然而，CPU 缓存无法明确管理，而共享内存可以。共享内存的延迟比全局内存低一个数量级，带宽比全局内存高一个数量级。但共享内存的关键用途来自于块内线程可以共享内存访问。CUDA 程序员可以使用共享变量来保存在内核执行阶段中多次重复使用的数据。此外，由于同一块内的线程可以共享结果，这有助于避免冗余计算。直到 9.0 版本，CUDA Toolkit 没有提供可靠的通信机制来在不同块的线程之间进行通信。我们将在后续章节中更详细地介绍 CUDA 9.0 通信机制。目前，我们将假设在 CUDA 中只能通过使用共享内存来实现线程之间的通信。

共享内存上的矩阵转置

用于演示共享内存的最原始的例子之一是矩阵转置。矩阵转置是一个内存绑定的操作。以下代码片段使用matrix_transpose_naive内核，展示了矩阵转置内核的示例实现：

__global__ void matrix_transpose_naive(int *input, int *output) {
     int indexX = threadIdx.x + blockIdx.x * blockDim.x;
     int indexY = threadIdx.y + blockIdx.y * blockDim.y;
     int index = indexY * N + indexX;
     int transposedIndex = indexX * N + indexY;
     output[index] = input[transposedIndex];
}

上述代码展示了使用全局内存的矩阵转置的朴素实现。如果以朴素的方式实现，这将导致在读取矩阵或写入矩阵时出现未协调的访问。在 V100 PCIe 16 GB 卡上，内核的执行时间约为 60 微秒。

根据以下步骤配置您的环境：

准备您的 GPU 应用程序。此代码可以在02_memory_overview/02_matrix_transpose中找到。
使用nvcc编译器编译您的应用程序，然后使用nvprof编译器对其进行分析。以下命令是对此的nvcc命令的一个示例。然后，我们使用nvprof命令对应用程序进行分析。还传递了--analysis-metrics标志以获取内核的指标。
生成的配置文件，即matrix_transpose.prof，然后加载到 NVIDIA Visual Profiler 中。用户需要从“文件|打开”菜单中加载分析输出。还要记得选择“所有文件”作为文件名选项的一部分：

$ nvcc -o matrix_transpose ./matrix_transpose.cu
$ nvcc -o conflict_solved ./conflict_solved.cu
$ nvprof --analysis-metrics --export-profile matrix_transpose.prof ./matrix_transpose
$ nvprof --analysis-metrics --export-profile conflict_solved.prof ./conflict_solved

以下截图显示了性能分析的输出。输出清楚地表明对全局内存的访问是未协调的，这是需要解决的关键指标，以便我们可以提高性能：

解决这个问题的一种方法是利用高带宽和低延迟的内存，比如共享内存。这里的诀窍是以协调的方式从全局内存读取和写入。在这里，对共享内存的读取或写入可以是未协调的模式。使用共享内存会带来更好的性能，时间缩短到 21 微秒，这是 3 倍的加速时间：

__global__ void matrix_transpose_shared(int *input, int *output) {

    __shared__ int sharedMemory [BLOCK_SIZE] [BLOCK_SIZE];

    //global index
     int indexX = threadIdx.x + blockIdx.x * blockDim.x;
     int indexY = threadIdx.y + blockIdx.y * blockDim.y;

    //transposed global memory index
     int tindexX = threadIdx.x + blockIdx.y * blockDim.x;
     int tindexY = threadIdx.y + blockIdx.x * blockDim.y;

    //local index
     int localIndexX = threadIdx.x;
     int localIndexY = threadIdx.y;
     int index = indexY * N + indexX;
     int transposedIndex = tindexY * N + tindexX;

    //transposed the matrix in shared memory. 
    // Global memory is read in coalesced fashion
     sharedMemory[localIndexX][localIndexY] = input[index];
     __syncthreads();

    //output written in global memory in coalesed fashion.
     output[transposedIndex] = sharedMemory[localIndexY][localIndexX];
}

上述代码片段显示了使用共享内存的矩阵转置的实现。全局内存读取/写入是协调的，而转置发生在共享内存中。

银行冲突及其对共享内存的影响

与使用全局内存相比的良好加速并不一定意味着我们有效地使用了共享内存。如果我们转换从引导分析到未引导分析的分析器输出，即matrix_transpose.prof，我们将看到共享内存访问模式显示出对齐问题，如下截图所示：

我们可以看到分析器显示了共享内存的非最佳使用，这是银行冲突的一个迹象。

为了有效地理解这个对齐问题，重要的是要理解bank的概念。共享内存被组织成 bank 以实现更高的带宽。每个 bank 可以在一个周期内服务一个地址。内存可以为它有的 bank 提供多个同时访问。Volta GPU 有 32 个 bank，每个 bank 宽度为 4 字节。当一个数组存储在共享内存中时，相邻的 4 字节单词会进入连续的 bank，如下图所示：

上述图中的逻辑视图显示了数据在共享内存中的存储方式。

warp 内的线程对 bank 的多个同时访问会导致 bank 冲突。换句话说，当 warp 内的两个或多个线程访问同一个 bank 中的不同 4 字节单词时，就会发生 bank 冲突。从逻辑上讲，这是当两个或多个线程访问同一个 bank 中的不同行时。以下图示例展示了不同n-way bank 冲突的例子。最坏的情况是 32-way 冲突 | 31 次重播 - 每次重播都会增加一些延迟：

上述情景显示了来自同一个 warp 的线程访问驻留在不同 bank 中的相邻 4 字节元素，导致没有 bank 冲突。看一下下图：

这是另一个没有银行冲突的场景，同一 warp 的线程访问随机的 4 字节元素，这些元素位于不同的银行中，因此没有银行冲突。由于共享内存中的 2 路银行冲突，顺序访问如下图所示：

前面的图表显示了一个场景，其中来自同一 warp 的线程 T0 和 T1 访问同一银行中的 4 字节元素，因此导致了 2 路银行冲突。

在前面的矩阵转置示例中，我们利用了共享内存来获得更好的性能。然而，我们可以看到 32 路银行冲突。为了解决这个问题，可以使用一种称为填充的简单技术。所有这些都是在共享内存中填充一个虚拟的，即一个额外的列，这样线程就可以访问不同的银行，从而获得更好的性能：

__global__ void matrix_transpose_shared(int *input, int *output) {

     __shared__ int sharedMemory [BLOCK_SIZE] [BLOCK_SIZE + 1];

    //global index
     int indexX = threadIdx.x + blockIdx.x * blockDim.x;
     int indexY = threadIdx.y + blockIdx.y * blockDim.y;

    //transposed index
     int tindexX = threadIdx.x + blockIdx.y * blockDim.x;
     int tindexY = threadIdx.y + blockIdx.x * blockDim.y;
     int localIndexX = threadIdx.x;
     int localIndexY = threadIdx.y;
     int index = indexY * N + indexX;
     int transposedIndex = tindexY * N + tindexX;

    //reading from global memory in coalesed manner 
    // and performing tanspose in shared memory
     sharedMemory[localIndexX][localIndexY] = input[index];

    __syncthreads();

    //writing into global memory in coalesed fashion 
    // via transposed data in shared memory
     output[transposedIndex] = sharedMemory[localIndexY][localIndexX];
}

前面的代码片段中，我们使用了matrix_transpose_shared内核，展示了填充的概念，这样可以消除银行冲突，从而更好地利用共享内存带宽。像往常一样，运行代码并借助可视化分析器验证这种行为。通过这些更改，您应该看到内核的时间减少到 13 微秒，这进一步提高了 60%的速度。

在本节中，我们看到了如何最大限度地利用共享内存，它提供了读写访问作为一个临时存储器。但有时，数据只是只读输入，不需要写访问。在这种情况下，GPU 提供了一种称为纹理内存的最佳内存。我们将在下一章中详细介绍这一点，以及它为开发人员提供的其他优势。我们将在下一节中介绍只读数据。

只读数据/缓存

根据内存名称，只读缓存适合存储只读数据，并且在内核执行过程中不会发生更改。该缓存针对此目的进行了优化，并且根据 GPU 架构，释放并减少了其他缓存的负载，从而提高了性能。在本节中，我们将详细介绍如何利用只读缓存，以及如何使用图像处理代码示例进行图像调整。

GPU 中的所有线程都可以看到只读数据。对于 GPU 来说，这些数据被标记为只读，这意味着对这些数据的任何更改都会导致内核中的未指定行为。另一方面，CPU 对这些数据具有读写访问权限。

传统上，这个缓存也被称为纹理缓存。虽然用户可以显式调用纹理 API 来利用只读缓存，但是在最新的 GPU 架构中，开发人员可以在不显式使用 CUDA 纹理 API 的情况下利用这个缓存。使用最新的 CUDA 版本和像 Volta 这样的 GPU，标记为const __restrict__的内核指针参数被视为只读数据，通过只读缓存数据路径传输。开发人员还可以通过__ldg内在函数强制加载这个缓存。

只读数据在算法要求整个 warp 读取相同地址/数据时理想地使用，这主要导致每个时钟周期对所有请求数据的线程进行广播。纹理缓存针对 2D 和 3D 局部性进行了优化。随着线程成为同一 warp 的一部分，从具有 2D 和 3D 局部性的纹理地址读取数据往往会获得更好的性能。纹理在要求随机内存访问的应用程序中已被证明是有用的，特别是在 Volta 架构之前的显卡中。

纹理支持双线性和三线性插值，这对于图像处理算法如缩放图像特别有用。

下图显示了一个 warp 内的线程访问空间位置在 2D 空间中的元素的示例。纹理适用于这类工作负载：

现在，让我们看一个关于缩放的小型实际算法，以演示纹理内存的使用。

计算机视觉-使用纹理内存进行图像缩放

我们将使用图像缩放作为示例来演示纹理内存的使用。图像缩放的示例如下截图所示：

图像缩放需要在 2 维中插值图像像素。纹理提供了这两个功能（插值和对 2D 局部性的高效访问），如果直接通过全局内存访问，将导致内存访问不连续。

根据以下步骤配置您的环境：

准备您的 GPU 应用程序。此代码可以在02_memory_overview/03_image_scaling中找到。
使用nvcc编译器编译您的应用程序，使用以下命令：

$nvcc -c scrImagePgmPpmPackage.cpp 
$nvcc -c image_scaling.cu
$nvcc -o image_scaling image_scaling.o scrImagePgmPpmPackage.o

scrImagePgmPpmPackage.cpp文件包含了读取和写入.pgm扩展名图像的源代码。纹理代码位于image_scaling.cu中。

用户可以使用 IrfanView（www.irfanview.com/main_download_engl.htm）等查看器来查看pgm文件，这些查看器是免费使用的。

主要有四个步骤是必需的，以便我们可以使用纹理内存：

声明纹理内存。
将纹理内存绑定到纹理引用。
在 CUDA 内核中使用纹理引用读取纹理内存。
从纹理引用中解绑纹理内存。

以下代码片段显示了我们可以使用的四个步骤来使用纹理内存。从 Kepler GPU 架构和 CUDA 5.0 开始，引入了一项名为无绑定纹理的新功能。这暴露了纹理对象，它基本上是一个可以传递给 CUDA 内核的 C++对象。它们被称为无绑定，因为它们不需要手动绑定/解绑，这是早期 GPU 和 CUDA 版本的情况。纹理对象使用cudaTextureObject_t类 API 声明。现在让我们通过这些步骤：

首先，声明纹理内存：

texture<unsigned char, 2, cudaReadModeElementType> tex;

创建一个通道描述，我们在链接到纹理时将使用：

cudaArray* cu_array;
cudaChannelFormatKind kind = cudaChannelFormatKindUnsigned;
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(8, 0, 0, 0, kind);

然后，指定纹理对象参数：

struct cudaTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc)); 
//set the memory to zero
texDesc.addressMode[0] = cudaAddressModeClamp; 
// setting the x dimension addressmode to Clamp
texDesc.addressMode[1] = cudaAddressModeClamp; 
//Setting y dimension addressmode to Clamp
texDesc.filterMode = cudaFilterModePoint; 
// Filter mode set to Point
texDesc.readMode = cudaReadModeElementType; 
// Reading element type and not interpolated
texDesc.normalizedCoords = 0;

接下来，在 CUDA 内核中从纹理引用中读取纹理内存：

imageScaledData[index] = tex2D<unsigned char>(texObj,(float)(tidX*scale_factor),(float)(tidY*scale_factor));

最后，销毁纹理对象：

cudaDestroyTextureObject(texObj);

纹理内存的重要方面，它们像配置一样由开发人员设置，如下所示：

纹理维度：这定义了纹理是作为 1D、2D 还是 3D 数组寻址。纹理中的元素也称为纹素。深度、宽度和高度也被设置以定义每个维度。请注意，每个 GPU 架构都定义了可接受的每个维度的最大尺寸。
纹理类型：这定义了基本整数或浮点纹素的大小。
纹理读取模式：纹理的读取模式定义了元素的读取方式。它们可以以NormalizedFloat或ModeElement格式读取。标准化浮点模式期望在[0.0 1.0]和[-1.0 1.0]范围内的索引，对于无符号整数和有符号整数类型。
纹理寻址模式：纹理的一个独特特性是它如何寻址超出范围的访问。这听起来可能很不寻常，但实际上在许多图像算法中非常常见。例如，如果您正在通过平均相邻像素来应用插值，那么边界像素的行为应该是什么？纹理为开发人员提供了这个选项，以便他们可以选择将超出范围视为夹紧、包裹或镜像。在调整大小的示例中，我们已将其设置为夹紧模式，这基本上意味着超出范围的访问被夹紧到边界。
纹理过滤模式：设置模式定义在获取纹理时如何计算返回值。支持两种类型的过滤模式：cudaFilterModePoint和cudaFilterModeLinear。当设置为线性模式时，可以进行插值（1D 的简单线性，2D 的双线性和 3D 的三线性）。仅当返回类型为浮点类型时，线性模式才有效。另一方面，ModePoint不执行插值，而是返回最近坐标的纹素。

在本节中引入纹理内存的关键目的是为您提供其用法的示例，并向您展示纹理内存的有用之处。它提供了不同配置参数的良好概述。有关更多信息，请参阅 CUDA API 指南（docs.nvidia.com/cuda/cuda-runtime-api/index.html）。

在本节中，我们通过示例描述了使用纹理内存的目的。在下一节中，我们将研究 GPU 内存中最快（最低延迟）的可用内存（寄存器）。与 CPU 相比，GPU 中富裕地存在这种内存。

GPU 中的寄存器

CPU 和 GPU 架构之间的一个基本区别是 GPU 中寄存器的丰富性相对于 CPU。这有助于线程将大部分数据保存在寄存器中，从而减少上下文切换的延迟。因此，使这种内存达到最佳状态也很重要。

寄存器的范围是单个线程。在 GRID 中为所有启动的线程创建变量的私有副本。每个线程都可以访问其变量的私有副本，而其他线程的私有变量则无法访问。例如，如果使用 1,000 个线程启动内核，那么作为线程范围的变量将获得其自己的变量副本。

作为内核的一部分声明的局部变量存储在寄存器中。中间值也存储在寄存器中。每个 SM 都有一组固定的寄存器。在编译期间，编译器（nvcc）尝试找到每个线程的最佳寄存器数量。如果寄存器数量不足，通常发生在 CUDA 内核较大且具有许多局部变量和中间计算时，数据将被推送到本地内存，该内存可以位于 L1/L2 缓存中，甚至更低的内存层次结构中，例如全局内存。这也被称为寄存器溢出。每个线程的寄存器数量在 SM 上可以激活多少个块和线程方面起着重要作用。这个概念在下一章中有详细介绍，该章节专门讨论了占用率。一般来说，建议不要声明大量不必要的局部变量。如果寄存器限制了可以在 SM 上调度的线程数量，那么开发人员应该考虑通过将内核拆分为两个或更多个（如果可能）来重新构建代码。

作为vecAdd内核的一部分声明的变量存储在寄存器内存中。传递给内核的参数，即A、B和C，指向全局内存，但变量本身存储在基于 GPU 架构的共享内存或寄存器中。以下图显示了 CUDA 内存层次结构和不同变量类型的默认位置：

到目前为止，我们已经看到了关键内存层次结构（全局、纹理、共享和寄存器）的目的和最佳用法。在下一节中，我们将看一些可以提高应用程序性能并增加开发人员在编写 CUDA 程序时的生产力的 GPU 内存的优化和特性。

固定内存

现在是时候回想一下数据的路径，即从 CPU 内存到 GPU 寄存器，最终由 GPU 核心用于计算。尽管 GPU 具有更高的计算性能和更高的内存带宽，但由于 CPU 内存和 GPU 内存之间的传输，应用程序获得的加速效果可能会变得规范化。数据传输是通过总线/链接/协议进行的，例如 PCIe（对于英特尔和 AMD 等 CPU 架构）或 NVLink（对于 OpenPower Foundation 的power等 CPU 架构）。

为了克服这些瓶颈，建议采用以下技巧/指南：

首先，建议在可能的情况下尽量减少主机和设备之间传输的数据量。这甚至可能意味着将顺序代码的一部分作为 GPU 上的内核运行，与在主机 CPU 上顺序运行相比，几乎没有或没有加速。
其次，通过利用固定内存，重要的是在主机和设备之间实现更高的带宽。
建议将小的传输批量成一个大的传输。这有助于减少调用数据传输 CUDA API 所涉及的延迟，根据系统配置，这可能从几微秒到几毫秒不等。
最后，应用程序可以利用异步数据传输来重叠内核执行和数据传输。

我们将在本节中更详细地介绍固定内存传输。异步传输将在第四章中更详细地介绍，内核执行模型和优化策略，在那里我们将使用一个称为 CUDA 流的概念。

带宽测试-固定与分页

默认情况下，称为malloc()的内存分配 API 分配的是可分页的内存类型。这意味着，如果需要，作为页面映射的内存可以被其他应用程序或操作系统本身交换出去。因此，大多数设备，包括 GPU 和其他设备（如 InfiniBand 等），也位于 PCIe 总线上，都希望在传输之前将内存固定。默认情况下，GPU 将不访问可分页内存。因此，当调用内存传输时，CUDA 驱动程序会分配临时固定内存，将数据从默认可分页内存复制到此临时固定内存，然后通过设备内存控制器（DMA）将其传输到设备。

这个额外的步骤不仅会增加延迟，还有可能会将请求的页面传输到已经被交换并需要重新传输到 GPU 内存的 GPU 内存。

为了了解使用固定内存的影响，让我们尝试编译和运行一段示例代码。这已作为 CUDA 示例的一部分提供。根据以下步骤配置您的环境：

准备您的 GPU 应用程序。此代码位于<CUDA_SAMPLES_DIR>/1_Utilities/bandwidthTest中。
使用make命令编译您的应用程序。
以分页和固定两种模式运行可执行文件，如下所示：

$make
$./bandwidthTest --mode=shmoo --csv --memory=pageable > pageable.csv
$./bandwidthTest --mode=shmoo --csv --memory=pinned >  pinned.csv

注意，CUDA_SAMPLES_DIR是 CUDA 安装所在目录的路径。

正如我们所看到的，与之前的代码相比，关键的变化是我们迄今为止编写的是数据分配 API。以下代码片段显示了使用cudaMallocHost API 而不是malloc来分配内存：

cudaError_t status = cudaMallocHost((void**)&h_aPinned, bytes);
if (status != cudaSuccess)
 printf("Error allocating pinned host memory\n");

cudaMallocHost API 使内存成为固定内存而不是分页内存。虽然分配 API 已更改，但我们仍然可以使用相同的数据传输 API，即cudaMemcpy()。现在，重要的问题是，固定内存是什么，为什么它提供更好的带宽？我们将在下一节中介绍这个问题。

性能的影响可以从带宽测试的输出中看出。我们已经将结果绘制成图表，以便您可以轻松理解影响。x轴显示了以 KB 为单位传输的数据，而y轴显示了以 MB/sec 为单位的实现带宽。

第一个图是主机到设备的传输，而第二个图是设备到主机的传输。您将看到的第一件事是可实现的最大带宽约为 12 GB/sec。PCIe Gen3 的理论带宽为 16 GB/sec，但实际可实现的范围在 12 GB/sec 左右。可实现的带宽高度取决于系统（主板、CPU、PCIe 拓扑等）：

如您所见，对于固定内存，在较小的传输大小时带宽始终更高，而在可分页内存中，随着数据大小的增加，带宽变得相等，因为驱动程序和 DMA 引擎开始通过应用诸如重叠的概念来优化传输。尽管建议使用固定内存，但过度使用也有缺点。为应用程序分配整个系统内存作为固定内存可能会降低整体系统性能。这是因为它会占用其他应用程序和操作系统任务可用的页面。应该固定的正确大小非常依赖于应用程序和系统，并且没有通用的公式可用。我们能做的最好的事情是在可用系统上测试应用程序并选择最佳的性能参数。

此外，重要的是要了解新的互连技术，如 NVLink，为受这些数据传输限制的应用程序提供了更高的带宽和更低的延迟。目前，CPU 和 GPU 之间的 NVLink 仅与 Power CPU 一起提供。

在本节中，我们将看看如何提高 CPU 和 GPU 之间的数据传输速度。现在我们将继续利用 CUDA 的一个新特性，称为统一内存，这有助于提高编写 CUDA 程序的开发人员的生产力。

统一内存

随着每一次新的 CUDA 和 GPU 架构发布，都会添加新的功能。这些新功能提供了更高的性能和更便捷的编程，或者允许开发人员实现新的算法，否则无法使用 CUDA 在 GPU 上进行移植。从 CUDA 6.0 开始发布的一个重要功能是统一内存，从 Kepler GPU 架构开始实现。在本章中，我们将统一内存称为 UM。

以更简单的话来说，UM 为用户提供了一个单一内存空间的视图，所有 GPU 和 CPU 都可以访问该空间。下图对此进行了说明：

在本节中，我们将介绍如何使用 UM，优化它，并突出利用它的关键优势。与全局内存访问一样，如果以不连续的方式进行，会导致性能不佳，如果未正确使用 UM 功能，也会导致应用程序整体性能下降。我们将采取逐步的方法，从一个简单的程序开始，并在此基础上构建，以便我们可以理解 UM 及其对性能的影响。

让我们尝试编译和运行一些示例代码。根据以下步骤配置您的环境：

准备您的 GPU 应用程序。此代码可以在02_memory_overview/unified_memory中找到。
使用以下nvcc命令编译您的应用程序：

$nvcc -o unified_simple.out unified_memory.cu
$nvcc -o unified_initialized.out unified_memory_initialized.cu
$nvcc -o unified_prefetch.out unified_memory_prefetch.cu
$nvcc -o unified_64align.out unified_memory_64align.cu

请注意，本节中显示的结果是针对 Tesla P100 卡的。当在其他架构（如 Kepler）上运行相同的代码时，预计会产生不同的结果。本节的重点是最新的架构，如 Pascal 和 Volta。

了解统一内存页面分配和传输

让我们从 UM 的朴素实现开始。代码的第一部分unified_memory.cu演示了这个概念的基本用法。代码中的关键更改是使用cudaMallocManaged()API 来分配内存，而不是使用malloc，如下面的代码片段所示：

float *x, *y;
int size = N * sizeof(float);
...
cudaMallocManaged(&x, size);
cudaMallocManaged(&y, size);
...

 for (int ix = 0; ix < N; ix++) {
    x[ix] = rand()%10;
    y[ix] = rand()%20;
  }
...

 add<<<numBlocks, blockSize>>>(x, y, N);

如果我们仔细查看源代码，我们会发现x和y变量只被分配一次并指向统一内存。同一个指针被发送到 GPU 的add<<<>>>()内核，并且在 CPU 中使用for循环进行初始化。这对程序员来说非常简单，因为他们不需要跟踪指针是指向 CPU 内存还是 GPU 内存。但这是否意味着我们能获得良好的性能或传输速度呢？不一定，所以让我们尝试通过对这段代码进行性能分析来深入了解，如下面的屏幕截图所示：

我们使用以下命令来获取性能分析输出：

$ nvprof ./unified_simple.out

正如预期的那样，大部分时间都花在了add<<<>>>内核上。让我们尝试理论计算带宽。我们将使用以下公式来计算带宽：

带宽 = 字节/秒 = (3 * 4,194,304 字节 * 1e-9 字节/GB) / 2.6205e-3 秒 = 5 GB/s

如您所见，P100 提供了 720 GB/s 的理论带宽，而我们只能实现 5 GB/s，这实在是太差了。您可能想知道为什么我们只计算内存带宽。这是因为应用程序受内存限制，因为它完成了三次内存操作和仅一次加法。因此，只集中在这个方面是有意义的。

从 Pascal 卡开始，cudaMallocManaged()不再分配物理内存，而是基于首次触摸的基础上分配内存。如果 GPU 首次触摸变量，页面将被分配并映射到 GPU 页表；否则，如果 CPU 首次触摸变量，它将被分配并映射到 CPU。在我们的代码中，x和y变量在 CPU 中用于初始化。因此，页面被分配给 CPU。在add<<<>>>内核中，当访问这些变量时，会发生页面错误，并且页面迁移的时间被添加到内核时间中。这是内核时间高的根本原因。现在，让我们深入了解页面迁移的步骤。

页面迁移中完成的操作顺序如下：

首先，我们需要在 GPU 和 CPU 上分配新页面（首次触摸）。如果页面不存在并且映射到另一个页面，会发生设备页表页错误。当在 GPU 中访问当前映射到 CPU 内存的page 2中的x时，会发生页面错误。请看下图：

接下来，CPU 上的旧页面被取消映射，如下图所示：

接下来，数据从 CPU 复制到 GPU，如下图所示：

最后，新页面在 GPU 上映射，旧页面在 CPU 上释放，如下图所示：

GPU 中的转换后备缓冲器（TLB）与 CPU 中的类似，执行从物理地址到虚拟地址的地址转换。当发生页面错误时，相应 SM 的 TLB 被锁定。这基本上意味着新指令将被暂停，直到执行前面的步骤并最终解锁 TLB。这是为了保持一致性并在 SM 内维护内存视图的一致状态。驱动程序负责删除这些重复项，更新映射并传输页面数据。正如我们之前提到的，所有这些时间都被添加到总体内核时间中。

所以，我们现在知道问题所在。但解决方案是什么呢？为了解决这个问题，我们将采用两种方法：

首先，我们将在 GPU 上创建一个初始化内核，以便在add<<<>>>内核运行期间没有页面错误。然后，我们将通过利用每页的 warp 概念来优化页面错误。
我们将预取数据。

我们将在接下来的部分中介绍这些方法。

使用每页 warp 优化统一内存

让我们从第一种方法开始，即初始化内核。如果你看一下unified_memory_initialized.cu文件中的源代码，我们在那里添加了一个名为init<<<>>>的新内核，如下所示：

__global__ void init(int n, float *x, float *y) {
 int index = threadIdx.x + blockIdx.x * blockDim.x;
 int stride = blockDim.x * gridDim.x;
 for (int i = index; i < n; i += stride) {
   x[i] = 1.0f;
   y[i] = 2.0f;
  }
}

通过在 GPU 本身添加一个初始化数组的内核，页面在init<<<>>>内核中首次被触摸时被分配和映射到 GPU 内存。让我们来看看这段代码的性能分析结果输出，其中显示了初始化内核的性能分析输出：

我们使用以下命令获取了性能分析输出

nvprof ./unified_initialized.out

正如你所看到的，add<<<>>>内核的时间减少到了 18 微秒。这有效地给了我们以下内核带宽：

带宽 = 字节/秒 = (3 * 4,194,304 字节 * 1e-9 字节/GB) / 18.84e-6 秒 = 670 GB/s

这个带宽是你在非统一内存场景中所期望的。正如我们从前面截图中的天真实现中所看到的，性能分析输出中没有主机到设备的行。然而，你可能已经注意到，即使add<<<>>>内核的时间已经减少，init<<<>>>内核也没有成为占用最长时间的热点。这是因为我们在init<<<>>>内核中首次触摸内存。此外，你可能想知道这些 GPU 错误组是什么。正如我们之前讨论的，个别页面错误可能会根据启发式规则和访问模式进行分组，以提高带宽。为了进一步深入了解这一点，让我们使用--print-gpu-trace重新对代码进行分析，以便我们可以看到个别页面错误。正如你从以下截图中所看到的，GPU 跟踪显示了错误的整体跟踪和发生错误的虚拟地址：

我们使用以下命令获取了性能分析输出：

$ nvprof --print-gpu-trace ./unified_initialized.out

第二行显示了相同页面的 11 个页面错误。正如我们之前讨论的，驱动程序的作用是过滤这些重复的错误并只传输每个页面一次。在复杂的访问模式中，通常驱动程序没有足够的信息来确定哪些数据可以迁移到 GPU。为了改善这种情况，我们将进一步实现每页 warp 的概念，这基本上意味着每个 warp 将访问位于相同页面中的元素。这需要开发人员额外的努力。让我们重新实现init<<<>>>内核。你可以在之前编译的unified_memory_64align.cu文件中看到这个实现。以下是内核的快照：

#define STRIDE_64K 65536
__global__ void init(int n, float *x, float *y) {
  int lane_id = threadIdx.x & 31;
  size_t warp_id = (threadIdx.x + blockIdx.x * blockDim.x) >> 5;
  size_t warps_per_grid = (blockDim.x * gridDim.x) >> 5;
  size_t warp_total = ((sizeof(float)*n) + STRIDE_64K-1) / STRIDE_64K;
  for(; warp_id < warp_total; warp_id += warps_per_grid) {
    #pragma unroll
    for(int rep = 0; rep < STRIDE_64K/sizeof(float)/32; rep++) {
      size_t ind = warp_id * STRIDE_64K/sizeof(float) + rep * 32 + lane_id;
      if (ind < n) {
        x[ind] = 1.0f;
        y[ind] = 2.0f;
      }
    }
  }
}

该内核显示索引是基于warp_id。 GPU 中的 warp 大小为 32，负责填充具有 64KB 范围的索引中的x和y变量，也就是说，warp 1 负责前 64KB 的部分，而 warp 2 负责接下来 64KB 的元素。warp 中的每个线程循环（最内层的for循环）以填充相同 64KB 内的索引。让我们来看看这段代码的性能分析结果。正如我们从以下截图中的性能分析输出中所看到的，init<<<>>>内核的时间已经减少，GPU 错误组也大大减少：

我们可以通过使用--print-gpu-trace运行分析器来重新确认这一点：

$ nvprof --print-gpu-trace ./unified_64align.out

以下截图清楚地显示了 GPU 每页的页面错误已经减少：

统一内存的优化使用数据预取

现在，让我们看一个更简单的方法，称为数据预取。CUDA 的一个关键特点是它为开发人员提供了不同的方法，从最简单的方法到需要忍者编程技能的方法。数据预取基本上是对驱动程序的提示，以在使用设备之前预取我们认为将在设备中使用的数据。CUDA 为此目的提供了一个名为cudaMemPrefetchAsync()的预取 API。要查看其实现，请查看我们之前编译的unified_memory_prefetch.cu文件。以下代码片段显示了此代码的快照：

// Allocate Unified Memory -- accessible from CPU or GPU
 cudaMallocManaged(&x, N*sizeof(float));  cudaMallocManaged(&y, N*sizeof(float));
// initialize x and y arrays on the host
 for (int i = 0; i < N; i++) {  x[i] = 1.0f;  y[i] = 2.0f;  } 
//prefetch the memory to GPU
cudaGetDevice(&device);
cudaMemPrefetchAsync(x, N*sizeof(float), device, NULL);
cudaMemPrefetchAsync(y, N*sizeof(float), device, NULL); 
...
 add<<<numBlocks, blockSize>>>(N, x, y);
//prefetch the memory to CPU
 cudaMemPrefetchAsync(y, N*sizeof(float), cudaCpuDeviceId, NULL);
 // Wait for GPU to finish before accessing on host
 cudaDeviceSynchronize();
...
for (int i = 0; i < N; i++)
 maxError = fmax(maxError, fabs(y[i]-3.0f));

代码非常简单，而且解释自己。概念相当简单：在已知将在特定设备上使用哪些内存的情况下，可以预取内存。让我们来看一下在以下截图中显示的分析结果。

正如我们所看到的，add<<<>>>内核提供了我们期望的带宽：

统一内存是一个不断发展的功能，随着每个 CUDA 版本和 GPU 架构的发布而发生变化。预计您通过访问最新的 CUDA 编程指南（docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#um-unified-memory-programming-hd）来保持自己的信息。

到目前为止，我们已经看到了 UM 概念的用处，它不仅提供了编程的便利（不需要使用 CUDA API 显式管理内存），而且在移植原本不可能移植到 GPU 上的应用程序或者原本移植困难的应用程序时更加强大和有用。使用 UM 的一个关键优势是超额订阅。与 CPU 内存相比，GPU 内存非常有限。最新的 GPU（Volta 卡 V100）每个 GPU 提供 32GB 的最大内存。借助 UM，多个 GPU 内存片段以及 CPU 内存可以被视为一个大内存。例如，拥有 16 个 Volta GPU 的 NVIDIA DGX2 机器，其内存大小为 323GB，可以被视为具有最大 512GB 大小的 GPU 内存集合。这对于诸如计算流体动力学（CFD）和分析等应用程序来说是巨大的优势。以前，很难将问题大小适应 GPU 内存，现在却是可能的。手动移动片段容易出错，并且需要调整内存大小。

此外，高速互连的出现，如 NVLink 和 NVSwitch，允许 GPU 之间进行高带宽和低延迟的快速传输。您实际上可以通过统一内存获得高性能！

数据预取，结合指定数据实际所在位置的提示，对于需要同时访问相同数据的多个处理器是有帮助的。在这种情况下使用的 API 名称是cudaMemAdvice()。因此，通过全面了解您的应用程序，您可以通过利用这些提示来优化访问。如果您希望覆盖某些驱动程序启发式方法，这些提示也是有用的。目前 API 正在采用的一些建议如下：

cudaMemAdviseSetReadMostly：顾名思义，这意味着数据大部分是只读的。驱动程序会创建数据的只读副本，从而减少页面错误。重要的是要注意，数据仍然可以被写入。在这种情况下，页面副本将变得无效，除了写入内存的设备：

// Sets the data readonly for the GPU
cudaMemAdvise(data, N, ..SetReadMostly, processorId); 
mykernel<<<..., s>>>(data, N);

cudaMemAdviseSetPreferredLocation：此建议将数据的首选位置设置为设备所属的内存。设置首选位置不会立即导致数据迁移到该位置。就像在以下代码中，mykernel<<<>>>将会出现页面错误并在 CPU 上生成数据的直接映射。驱动程序试图抵制将数据迁离设置的首选位置，使用cudaMemAdvise：

cudaMemAdvise(input, N, ..PreferredLocation, processorId); 
mykernel<<<..., s>>>(input, N);

cudaMemAdviseSetAccessedBy：这个建议意味着数据将被设备访问。设备将在 CPU 内存中创建输入的直接映射，不会产生页面错误：

cudaMemAdvise(input, N, ..SetAccessedBy, processorId); 
mykernel<<<..., s>>>(input, N);

在接下来的部分，我们将以整体的视角来看 GPU 中不同的内存是如何随着新的架构而发展的。

GPU 内存的演变

GPU 架构随着时间的推移发生了变化，内存架构也发生了相当大的变化。如果我们看一下过去四代的情况，会发现一些共同的模式，其中一些如下：

内存容量总体上已经提高了几个级别。
内存带宽和容量随着新一代架构的出现而增加。

以下表格显示了过去四代的属性：

内存类型	属性	Volta V100	Pascal P100	Maxwell M60	Kepler K80
寄存器	每个 SM 的大小	256 KB	256 KB	256 KB	256 KB
L1	大小	32...128 KiB	24 KiB	24 KiB	16...48 KiB
行大小	32	32 B	32 B	128 B
L2	大小	6144 KiB	4,096 KiB	2,048 KiB	1,536 Kib
行大小	64 B	32B	32B	32B
共享内存	每个 SMX 的大小	高达 96 KiB	64 KiB	64 KiB	48 KiB
每个 GPU 的大小	高达 7,689 KiB	3,584 KiB	1,536 KiB	624 KiB
理论带宽	13,800 GiB/s	9,519 GiB/s	2,410 GiB/s	2,912 GiB/s
全局内存	内存总线	HBM2	HBM2	GDDR5	GDDR5
大小	32,152 MiB	16,276 MiB	8,155 MiB	12,237 MiB
理论带宽	900 GiB/s	732 GiB/s	160 GiB/s	240 GiB/s

总的来说，前面的观察结果已经帮助 CUDA 应用在新的架构下运行得更快。但与此同时，CUDA 编程模型和内存架构也进行了一些根本性的改变，以便为 CUDA 程序员简化工作。我们观察到的一个这样的改变是纹理内存，之前在 CUDA 5.0 之前，开发人员必须手动绑定和解绑纹理，并且必须在全局声明。但在 CUDA 5.0 中，这是不必要的。它还取消了应用程序中开发人员可以拥有的纹理引用数量的限制。

我们还研究了 Volta 架构以及为简化开发人员编程而进行的一些根本性改变。Volta 的总容量是每个 SM 128 KB，比其上一代显卡 Pascal P100 多了七倍，这为开发人员提供了更大的缓存。此外，由于 Volta 架构中 L1 缓存的延迟要小得多，这使得它对频繁重用的数据具有高带宽和低延迟的访问。这样做的关键原因是让 L1 缓存操作获得共享内存性能的好处。共享内存的关键问题是需要开发人员显式控制。在使用 Volta 等新架构时，这种需求就不那么必要了。但这并不意味着共享内存变得多余。一些极客程序员仍然希望充分利用共享内存的性能，但许多其他应用程序不再需要这种专业知识。Pascal 和 Volta L1 缓存和共享内存之间的区别如下图所示：

前面的图表显示了与 Pascal 相比共享内存和 L1 缓存的统一。重要的是要理解，CUDA 编程模型从诞生以来几乎保持不变。尽管每个架构的内存容量、带宽或延迟都在变化，但相同的 CUDA 代码将在所有架构上运行。不过，随着这些架构变化，性能的影响肯定会发生变化。例如，在 Volta 之前利用共享内存的应用程序，与使用全局内存相比可能会看到性能提升，但在 Volta 中可能不会看到这样的加速，因为 L1 和共享内存的统一。

为什么 GPU 有缓存？

在这个演变过程中，还很重要的一点是要理解 CPU 和 GPU 缓存是非常不同的，而且有不同的用途。作为 CUDA 架构的一部分，我们通常在每个 SM 上启动数百到数千个线程。数万个线程共享 L2 缓存。因此，L1 和 L2 对每个线程来说都很小。例如，在每个 SM 上有 2,048 个线程，共有 80 个 SM，每个线程只能获得 64 字节的 L1 缓存和 38 字节的 L2 缓存。GPU 缓存中存储着许多线程访问的公共数据。这有时被称为空间局部性。一个典型的例子是当线程的访问是不对齐和不规则的时候。GPU 缓存可以帮助减少寄存器溢出和局部内存的影响，因为 CPU 缓存主要用于时间局部性。

总结

我们在本章开始时介绍了不同类型的 GPU 内存。我们详细讨论了全局、纹理和共享内存，以及寄存器。我们还看了 GPU 内存演变提供了哪些新功能，例如统一内存，这有助于提高程序员的生产力。我们看到了这些功能在最新的 GPU 架构（如 Pascal 和 Volta）中是如何实现的。

在下一章中，我们将深入讨论 CUDA 线程编程的细节，以及如何最优地启动不同的线程配置，以发挥 GPU 硬件的最佳性能。我们还将介绍新的 CUDA Toolkit 功能，例如用于灵活线程编程的协作组和 GPU 上的多精度编程。

第三章：CUDA 线程编程

CUDA 具有分层线程架构，因此我们可以控制 CUDA 线程的分组。了解它们在 GPU 上并行工作的方式有助于您编写并行编程代码并实现更好的性能。在本章中，我们将介绍 CUDA 线程操作及其与 GPU 资源的关系。作为实际经验，我们将研究并行减少算法，并看看如何通过使用优化策略来优化 CUDA 代码。

在本章中，您将学习 CUDA 线程在 GPU 中的操作：并行和并发线程执行，warp 执行，内存带宽问题，控制开销，SIMD 操作等等。

本章将涵盖以下主题：

层次化的 CUDA 线程操作
了解 CUDA 占用率
跨多个 CUDA 线程共享数据
识别应用程序的性能限制
最小化 CUDA warp 分歧效应
增加内存利用率和网格跨距循环
用于灵活线程处理的协作组
warp 同步编程
低/混合精度操作

技术要求

本章建议使用比 Pascal 架构更晚的 NVIDIA GPU 卡。换句话说，您的 GPU 的计算能力应等于或大于 60。如果您不确定您的 GPU 架构，请访问 NVIDIA 的 GPU 网站developer.nvidia.com/cuda-gpus，并确认您的 GPU 的计算能力。

在我们写这本书的时候，示例代码是使用 10.1 版本开发和测试的。一般来说，如果适用的话，建议使用最新的 CUDA 版本。

在本章中，我们将通过对代码进行性能分析来进行 CUDA 编程。如果您的 GPU 架构是图灵架构，建议安装 Nsight Compute 来对代码进行性能分析。它是免费的，您可以从developer.nvidia.com/nsight-compute下载。在我们写这本书的时候，这是性能分析工具的过渡时刻。您可以在第五章的使用 Nsight Compute 对内核进行性能分析部分了解其基本用法，CUDA 应用性能分析和调试。

CUDA 线程、块和 GPU

CUDA 编程中的基本工作单元是 CUDA 线程。基本的 CUDA 线程执行模型是单指令多线程（SIMT）。换句话说，内核函数的主体是单个 CUDA 线程的工作描述。但是，CUDA 架构执行具有相同操作的多个 CUDA 线程。

在概念上，多个 CUDA 线程以组的形式并行工作。CUDA 线程块是多个 CUDA 线程的集合。多个线程块同时运行。我们称线程块的组为网格。以下图表显示了它们之间的关系：

这些分层的 CUDA 线程操作与分层的 CUDA 架构相匹配。当我们启动 CUDA 内核时，每个流多处理器上会执行一个或多个 CUDA 线程块。此外，根据资源的可用性，一个流多处理器可以运行多个线程块。线程块中的线程数量和网格中的块数量也会有所不同。

流多处理器以任意和并发的方式执行线程块，执行尽可能多的 GPU 资源。因此，可并行执行的线程块数量取决于块需要的 GPU 资源量以及 GPU 资源的可用量。我们将在接下来的部分中介绍这一点。流多处理器的数量取决于 GPU 规格。例如，Tesla V100 为 80，RTX 2080（Ti）为 48。

CUDA 流多处理器以 32 个线程的组形式控制 CUDA 线程。一个组被称为warp。这样，一个或多个 warp 配置一个 CUDA 线程块。以下图显示了它们的关系：

小绿色框是 CUDA 线程，它们被 warp 分组。warp 是 GPU 架构的基本控制单元。因此，它的大小对 CUDA 编程具有隐式或显式的影响。例如，最佳线程块大小是在可以充分利用块的 warp 调度和操作的多个 warp 大小中确定的。我们称之为占用率，这将在下一节中详细介绍。此外，warp 中的 CUDA 线程并行工作，并具有同步操作。我们将在本章的Warp 级别基元编程部分讨论这一点。

利用 CUDA 块和 warp

现在，我们将研究 CUDA 线程调度及其使用 CUDA 的printf进行隐式同步。并行 CUDA 线程的执行和块操作是并发的。另一方面，从设备打印输出是一个顺序任务。因此，我们可以轻松地看到它们的执行顺序，因为对于并发任务来说输出是任意的，而对于并行任务来说是一致的。

我们将开始编写打印全局线程索引、线程块索引、warp 索引和 lane 索引的内核代码。为此，代码可以编写如下：

__global__ void index_print_kernel() {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int warp_idx = threadIdx.x / warpSize;
    int lane_idx = threadIdx.x & (warpSize - 1);

    if ((lane_idx & (warpSize/2 - 1)) == 0)
        //thread, block, warp, lane
        printf(" %5d\t%5d\t %2d\t%2d\n", idx, blockIdx.x, 
               warp_idx, lane_idx);
}

这段代码将帮助我们理解 warp 和 CUDA 线程调度的并发性。让我们让我们的代码从 shell 获取参数，以便轻松测试各种网格和线程块配置。

然后，我们将编写调用内核函数的主机代码：

int main() {
    int gridDim = 4, blockDim = 128;
    puts("thread, block, warp, lane");
    index_print_kernel<<< gridDim, blockDim >>>();
    cudaDeviceSynchronize();
}

最后，让我们编译代码，执行它，并查看结果：

nvcc -m64 -o cuda_thread_block cuda_thread_block.cu

以下结果是输出结果的一个示例。实际输出可能会有所不同：

$ ./cuda_thread_block.cu 4 128
thread, block, warp, lane
 64     0     2     0
 80     0     2    16
 96     0     3     0
 112     0     3    16
 0     0     0     0
 16     0     0    16
 ...
 352     2     3     0
 368     2     3    16
 288     2     1     0
 304     2     1    16

从结果中，您将看到 CUDA 线程以 warp 大小启动，并且顺序是不确定的。另一方面，lane 输出是有序的。从给定的结果中，我们可以确认以下事实：

无序块执行：第二列显示线程块的索引。结果表明，它不保证按照块索引的顺序执行。
无序 warp 索引与线程块：第三列显示块中 warp 的索引。warp 的顺序在块之间变化。因此，我们可以推断 warp 执行顺序没有保证。
在 warp 中执行的分组线程：第四列显示 warp 中的 lane。为了减少输出数量，应用程序限制只打印两个索引。从每个 warp 内的有序输出中，我们可以类比printf函数的输出顺序是固定的，因此没有倒置。

总之，CUDA 线程被分组为 32 个线程，它们的输出和 warp 的执行没有顺序。因此，程序员必须牢记这一点，以便进行 CUDA 内核开发。

理解 CUDA 占用率

CUDA 占用率是活动 CUDA warps 与每个流多处理器可以同时执行的最大 warps 的比率。一般来说，更高的占用率会导致更有效的 GPU 利用率，因为有更多的 warp 可用来隐藏停滞 warp 的延迟。然而，它也可能由于 CUDA 线程之间资源争用的增加而降低性能。因此，开发人员理解这种权衡是至关重要的。

找到最佳的 CUDA 占用率的目的是使 GPU 应用程序能够有效地使用 GPU 资源发出 warp 指令。GPU 在流多处理器上使用多个 warp 调度器调度多个 warp。当多个 warp 有效地调度时，GPU 可以隐藏 GPU 指令或内存延迟之间的延迟。然后，CUDA 核心可以执行连续从多个 warp 发出的指令，而未调度的 warp 必须等待，直到它们可以发出下一条指令。

开发人员可以使用两种方法确定 CUDA 占用率：

由 CUDA 占用率计算器确定的理论占用率：这个计算器是 CUDA 工具包提供的一个 Excel 表。我们可以从内核资源使用和 GPU 流多处理器理论上确定每个内核的占用率。
由 GPU 确定的实现占用率：实现占用率反映了在流多处理器上并发执行的 warp 的真实数量和最大可用 warp。这种占用率可以通过 NVIDIA 分析器进行度量分析来测量。

理论占用率可以被视为最大的上限占用率，因为占用率数字不考虑指令依赖性或内存带宽限制。

现在，让我们看看这个占用率和 CUDA C/C++之间的关系。

设置 NVCC 报告 GPU 资源使用

首先，我们将使用简单矩阵乘法（SGEMM）内核代码，如下所示：

__global__ void sgemm_gpu_kernel(const float *A, const float *B, 
        float *C, int N, int M, int K, alpha, float beta) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;

    float sum = 0.f;
    for (int i = 0; i < K; ++i) {
        sum += A[row * K + i] * B[i * K + col];
    }
    C[row * M + col] = alpha * sum + beta * C[row * M + col];
}

然后，我们将使用以下内核代码调用内核函数：

void sgemm_gpu(const float *A, const float *B, float *C,
            int N, int M, int K, float alpha, float beta) {
    dim3 dimBlock(BLOCK_DIM, BLOCK_DIM);
    dim3 dimGrid(M / dimBlock.x, N / dimBlock.y);
    sgemm_gpu_kernel<<< dimGrid, dimBlock >>>(A, B, C, N, M, K, alpha, beta);
}

您可能希望提供适当的 GPU 内存及其大小信息。我们将使用 2048 作为N，M和K。内存大小是该数字的平方。我们将把BLOCK_DIM设置为16。

现在，让我们看看如何使nvcc编译器报告内核函数的 GPU 资源使用情况。

Linux 设置

在 Linux 环境中，我们应该提供两个编译器选项，如下所示：

--resource-usage（--res-usage）：为 GPU 资源使用设置详细选项
-gencode：指定要编译和生成操作码的目标架构如下：
Turing：compute_75,sm_75
Volta：compute_70,sm_70
Pascal：compute_60,sm_60，compute_61,sm_61

如果您不确定您正在使用哪种架构，您可以从 CUDA GPU 网站上找到（developer.nvidia.com/cuda-gpus）。例如，nvcc编译命令可以有以下编译选项：

$ nvcc -m 64 --resource-usage \
 -gencode arch=compute_70,code=sm_70 \
 -I/usr/local/cuda/samples/common/inc \
 -o sgemm ./sgemm.cu

我们还可以编译代码以针对多个 GPU 架构，如下所示：

$ nvcc -m64 --resource-usage \
      -gencode arch=compute_70,code=sm_70 \
      -gencode arch=compute_75,code=sm_75 \
      -I/usr/local/cuda/samples/common/inc \
      -o sgemm ./sgemm.cu

如果您想使您的代码与新的 GPU 架构（Turing）兼容，您需要提供以下附加选项：

$ nvcc -m64 --resource-usage \
      -gencode arch=compute_70,code=sm_70 \
      -gencode arch=compute_75,code=sm_75 \
      -gencode arch=compute_75,code=compute_75 \
      -I/usr/local/cuda/samples/common/inc \
      -o sgemm ./sgemm.cu

如果您想了解更多关于这些选项的信息，您可以在这个文档中找到相关信息：docs.nvidia.com/cuda/turing-compatibility-guide/index.html#building-turing-compatible-apps-using-cuda-10-0。

现在，让我们编译源代码。我们可以从 NVCC 的输出中找到一个资源使用报告。以下结果是使用前面的命令生成的：

NVCC 为每个计算能力报告 CUDA 内核资源使用信息。在前面的输出截图中，我们可以看到每个线程的寄存器数量和常量内存使用情况。

Windows 设置

当我们开发 Windows 应用程序时，我们可以在 Visual Studio 项目属性对话框中设置这些设置。以下是该对话框的截图：

要打开此对话框，我们应该打开 debug_vs 属性页，然后在左侧面板上转到 CUDA C/C++ | 设备选项卡。然后，我们应该设置以下选项如下：

Verbose PTXAS Output: No | Yes
代码生成：更新选项以指定您的目标架构如下：
图灵：compute_75,sm_75
伏尔塔：compute_70,sm_70
帕斯卡：compute_60,sm_60;compute_61,sm_61

我们可以使用分号(;)指定多个目标架构。

现在，让我们构建源代码，我们将在 Visual Studio 的输出面板上看到 NVCC 的报告。然后，你会看到类似以下的输出：

这与 Linux 中 NVCC 的输出相同。

现在，让我们使用资源使用报告来分析内核的占用情况。

使用占用率计算器分析最佳占用率

实际上，我们可以使用 CUDA 占用率计算器，它是 CUDA 工具包提供的。使用这个，我们可以通过提供一些内核信息来获得理论上的占用率。计算器是一个 Excel 文件，你可以在以下位置找到它，根据你使用的操作系统：

Windows: C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\<cuda-version>\tools
Linux: /usr/local/cuda/tools
macOS: /Developer/NVIDIA/<cuda-version>/tools

以下是计算器的屏幕截图：

CUDA 占用率计算器

这个计算器有两部分：内核信息输入和占用信息输出。作为输入，它需要两种信息，如下所示：

GPU 的计算能力（绿色）
线程块资源信息（黄色）:
每个 CUDA 线程块的线程
每个 CUDA 线程的寄存器
每个块的共享内存

计算器在这里显示了 GPU 的占用信息：

GPU 占用数据（蓝色）
GPU 的 GPU 计算能力的物理限制（灰色）
每个块分配的资源（黄色）
每个流多处理器的最大线程块（黄色、橙色和红色）
根据三个关键的占用资源（线程、寄存器和每个块的共享内存），绘制占用限制图
图表上的红色三角形，显示当前的占用数据

现在，让我们把获得的信息放入计算器中。我们可以编辑 Excel 表格中的绿色和橙色区域：

输入你获得的内核资源信息，看看表格如何变化。

根据计算能力和输入数据，占用情况会发生变化，如下面的屏幕截图所示：

根据计算能力和输入数据的变化

蓝色区域显示了内核函数实现的占用率。在这个屏幕截图中，它显示了 100%的占用率。表格的右侧显示了 GPU 资源的占用率利用图：CUDA 线程、共享内存和寄存器。

一般来说，由于许多原因，内核代码不能达到 100%的理论占用率。然而，设置峰值占用率是有效利用 GPU 资源的开始。

占用率调整 - 限制寄存器使用

当内核算法复杂或处理数据类型为双精度时，CUDA 寄存器使用可能会增加。在这种情况下，由于活动 warp 大小有限，占用率会下降。在这种情况下，我们可以通过限制寄存器使用来增加理论上的占用率，并查看性能是否提高。

调整 GPU 资源使用的一种方法是在内核函数中使用__launch_bound__限定符。这告诉 NVCC 保证每个流多处理器的最大块大小的最小线程块。然后，NVCC 找到实现给定条件的最佳寄存器大小。如果你在编译时知道使你的算法有效运行的大小，你可以使用这个。标识符可以如下使用：

int maxThreadPerBlock = 256;
int minBlocksPerMultiprocessor = 2;
__global__ void
__launch_bound__ (maxThreadPerBlock, minBlocksPerMultiprocessor) foo_kernel() {
    ...
}

然后，编译器检查上限资源并减少每个块的限制资源使用。如果其资源使用没有超过上限，编译器会调整寄存器使用，如果 CUDA 可以调度额外的多处理器线程块，如果没有给出第二个参数。或者，编译器会增加寄存器使用以隐藏单线程指令延迟。

此外，我们可以简单地在应用程序级别限制占用寄存器的数量。--maxrregcount标志到NVCC将指定数量，编译器将重新排列寄存器使用。以下编译命令显示了如何在 Linux 终端中使用该标志：

$ nvcc -m64 -I/usr/local/cuda/samples/common/inc -gencode arch=compute_70,code=sm_70 --resource-usage --maxrregcount 24 -o sgemm ./sgemm.cu

但是，请记住，以这种方式限制寄存器使用可能会引入由寄存器限制引起的线程性能下降。即使编译器无法将其设置在限制之下，也可以将寄存器分割为本地内存，并且本地变量放置在全局内存中。

从分析器获取实现的占用

现在，我们可以使用 Visual Profiler 从分析的度量数据中获取实现的占用。单击目标内核时间轴条。然后，我们可以在属性面板中看到理论和实现的占用。我们还可以从内核延迟菜单中获取更多详细信息。以下屏幕截图显示了我们使用的示例代码的实现性能：

显示实现和理论占用的性能

通过这种占用调整，我们可以设计 CUDA 块大小，充分利用流多处理器中的 warp 调度。然而，这并没有解决我们在上一节中发现的 54.75%的内存限制问题。这意味着多处理器可能会停顿，无法掩盖由于受阻内存请求而产生的内存访问延迟。我们将在本章讨论如何优化这一点，并且在第七章《CUDA 中的并行编程模式》中，我们将讨论矩阵乘法优化。

理解并行归约

归约是一种简单但有用的算法，可以获得许多参数的公共参数。这个任务可以按顺序或并行完成。当涉及到并行处理到并行架构时，并行归约是获得直方图、均值或任何其他统计值的最快方式。

以下图表显示了顺序归约和并行归约之间的差异：

通过并行进行归约任务，可以将并行归约算法的总步骤减少到对数级别。现在，让我们开始在 GPU 上实现这个并行归约算法。首先，我们将使用全局内存实现一个简单的设计。然后，我们将使用共享内存实现另一个归约版本。通过比较这两种实现，我们将讨论是什么带来了性能差异。

使用全局内存的天真并行归约

归约的第一种基本方法是使用并行的 CUDA 线程，并使用全局内存共享归约输出。对于每次迭代，CUDA 内核通过将其大小减少两倍来从全局内存获取累积值。归约的工作如下图所示，显示了使用全局内存数据共享的天真并行归约：

这种方法在 CUDA 中很慢，因为它浪费了全局内存的带宽，并且没有利用任何更快的片上内存。为了获得更好的性能，建议使用共享内存来节省全局内存带宽并减少内存获取延迟。我们将讨论这种方法如何浪费带宽。

现在，让我们实现这个归约。首先，我们将编写归约内核函数，如下所示：

__global__ void naive_reduction_kernel
     (float *data_out, float *data_in, int stride, int size) {
     int idx_x = blockIdx.x * blockDim.x + threadIdx.x;
     if (idx_x + stride < size)
         data_out[idx_x] += data_in[idx_x + stride];
}

我们将在迭代过程中减半减小步长大小，直到stride大小为 1 时调用内核函数。

void naive_reduction(float *d_out, float *d_in, int n_threads, int size) {
    int n_blocks = (size + n_threads - 1) / n_threads;
    for (int stride = 1; stride < size; stride *= 2)
        naive_reduction_kernel<<<n_blocks, n_threads>>>(d_out, d_in, stride, size);
}

在这个实现中，内核代码使用跨距寻址获取设备内存并输出一个减少结果。主机代码触发每个步骤的减少内核，并且参数大小减半。我们不能有内部内核循环，因为 CUDA 不能保证线程块和流多处理器之间的同步操作。

使用共享内存减少内核

在这种减少中，每个 CUDA 线程块减少输入值，并且 CUDA 线程使用共享内存共享数据。为了进行适当的数据更新，它们使用块级内在同步函数__syncthreads()。然后，下一个迭代操作上一个减少结果。其设计如下图所示，显示了使用共享内存的并行减少：

黄点框表示 CUDA 线程块的操作范围。在这个设计中，每个 CUDA 线程块输出一个减少结果。

块级减少允许每个 CUDA 线程块进行减少，并输出单个减少输出。由于它不需要我们将中间结果保存在全局内存中，CUDA 内核可以将过渡值存储在共享内存中。这种设计有助于节省全局内存带宽并减少内存延迟。

与全局减少一样，我们将实现这个操作。首先，我们将编写内核函数，如下所示：

__global__ void reduction_kernel(float* d_out, float* d_in, 
                                 unsigned int size) {
    unsigned int idx_x = blockIdx.x * blockDim.x + threadIdx.x;

    extern __shared__ float s_data[];
    s_data[threadIdx.x] = (idx_x < size) ? d_in[idx_x] : 0.f;

    __syncthreads();

    // do reduction
    for (unsigned int stride = 1; stride < blockDim.x; stride *= 2) {
        // thread synchronous reduction
        if ( (idx_x % (stride * 2)) == 0 )
            s_data[threadIdx.x] += s_data[threadIdx.x + stride];

        __syncthreads();
    }

    if (threadIdx.x == 0)
        d_out[blockIdx.x] = s_data[0];
}

然后，我们将调用内核函数，如下所示：

void reduction(float *d_out, float *d_in, int n_threads, int size)
{
    cudaMemcpy(d_out, d_in, size * sizeof(float), cudaMemcpyDeviceToDevice);
    while(size > 1) {
        int n_blocks = (size + n_threads - 1) / n_threads;
        reduction_kernel
            <<< n_blocks, n_threads, n_threads * sizeof(float), 0 >>>
            (d_out, d_out, size);
        size = n_blocks;
    }
}

在这段代码中，我们提供了n_threads * sizeof (float)字节，因为每个 CUDA 线程将共享每个字节的单个变量。

编写性能测量代码

为了测量每个版本的性能，我们将使用 CUDA 示例timer辅助函数：

// Initialize timer
StopWatchInterface *timer;
sdkCreateTimer(&timer);
sdkStartTimer(&timer);

... Execution code ...

// Getting elapsed time
cudaDeviceSynchronize(); // Blocks the host until GPU finishes the work
sdkStopTimer(&timer);

// Getting execution time in micro-secondes
float execution_time_ms = sdkGetTimerValue(&timer)

// Termination of timer
sdkDeleteTimer(&timer);

这个函数集有助于在微秒级别测量执行时间。此外，建议在性能测量之前调用内核函数，以消除设备初始化开销。有关更详细的实现，请访问global_reduction.cu和reduction.cu文件中的实现代码。这些代码集在本章中用于评估优化效果以及分析器。

两种减少-全局和共享内存的性能比较

现在，我们可以比较两个并行减少操作的执行时间。性能可能会因 GPU 和实现环境而异。分别运行以下命令进行全局减少和使用共享内存进行减少：

# Reduction with global memory
$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o reduction_global ./reduction_global.cpp reduction_global_kernel.cu

# Reduction using shared memory
$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o reduction_shared ./reduction_shared.cpp reduction_shared_kernel.cu

使用我的 Tesla V100 PCIe 卡，两种减少的估计性能如下。元素数量为2²⁴个：

操作	估计时间（毫秒）	加速
原始方法（使用全局内存进行减少）	4.609	1.0x
使用共享内存的减少	0.624	7.4x

从这个结果中，我们可以看到在减少中使用共享内存共享数据如何快速返回输出。第一个实现版本在global_reduction.cu中，第二个版本在shared_reduction.cu中，所以您可以自行比较实现。

通过将减少与共享内存结合，我们可以显著提高性能。然而，我们无法确定这是否是我们可以获得的最大性能，并且不知道我们的应用程序有什么瓶颈。为了分析这一点，我们将在下一节中涵盖性能限制器。

识别应用程序的性能限制器

之前，我们看到了如何通过保存全局内存来使 CUDA 内核的性能受益。一般来说，使用片上缓存比使用片外内存更好。但是，我们无法确定这种简单类比是否还有很多优化空间。

性能限制因素显示了限制应用程序性能的因素，它最显著地限制了应用程序的性能。根据其分析信息，它分析了计算和内存带宽之间的性能限制因素。根据这些资源的利用率，应用程序可以被分类为四种类型：计算受限，带宽受限，延迟受限和计算和延迟受限。以下图表显示了这些类别与计算和内存利用率的关系：

在确定了限制因素之后，我们可以使用下一个优化策略。如果任一资源的利用率很高，我们可以专注于优化该资源。如果两者都未充分利用，我们可以从系统的 I/O 方面应用延迟优化。如果两者都很高，我们可以调查是否存在内存操作停顿问题和与计算相关的问题。

现在让我们看看如何获得利用率信息。

找到性能限制因素并进行优化

现在，让我们将此分析应用于两个减少实现。我们将对它们进行比较，并讨论共享内存如何有助于性能限制因素分析以改善性能。首先，让我们使用以下命令对基于全局内存的减少应用程序进行度量分析：

$ nvprof -o reduction_global.nvvp ./reduction_global 
$ nvprof --analysis-metrics -o reduction_global_metric.nvvp ./reduction_global

然后，我们将从 NVIDIA 分析器获得以下图表，显示了基于全局内存的第一个减少性能的限制因素：

在这张图表上，我们需要查看性能执行比来查看是否通过检查内核延迟分析来平衡。因为，如前图表所示，计算和内存之间的利用率差距很大，这可能意味着由于内存瓶颈，计算中会有很多延迟。以下图表显示了基于采样的分析结果，我们可以确定 CUDA 核心由于内存依赖而饥饿：

如您所见，由于内存等待，内核执行被延迟。现在，让我们基于共享内存对减少进行分析。我们可以使用以下命令来做到这一点：

$ nvprof -o reduction_shared.nvvp ./reduction_shared 
$ nvprof --analysis-metrics -o reduction_shared_metric.nvvp ./reduction_shared

然后，我们将获得以下图表，显示了基于共享内存的第二个减少性能的限制因素：

我们可以确定它是计算受限的，内存不会使 CUDA 核心饥饿。

现在让我们回顾我们的核心操作，以优化计算操作。以下代码显示了内核函数中的并行减少部分：

for (unsigned int stride = 1; stride < blockDim.x; stride *= 2) {
     if ( (idx_x % (stride * 2)) == 0 )
         s_data[threadIdx.x] += s_data[threadIdx.x + stride];
     __syncthreads();
 }

作为算术操作，模运算是一种重型操作。由于stride变量是2的指数倍数，因此可以用位操作替换，如下所示：

for (unsigned int stride = 1; stride < blockDim.x; stride *= 2) {
     if ( (idx_x & (stride * 2 - 1)) == 0 )
         s_data[threadIdx.x] += s_data[threadIdx.x + stride];
     __syncthreads();
 }

运行以下命令以查看优化后的输出：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o reduction_shared ./reduction_shared.cpp reduction_shared_kernel.cu

然后，新的估计时间为0.399 毫秒，我们可以实现更优化的性能，如下表所示：

操作	估计时间（毫秒）	加速比
原始方法（使用全局内存进行减少）	4.609	1.0 倍
使用共享内存进行减少	0.624	7.4 倍
将条件操作从`%`更改为`&`	0.399	11.55 倍

以下图表显示了更新后的性能限制因素：

我们可以确定其操作是计算和延迟受限。因此，我们可以确定我们可以通过优化计算效率来增加内存利用率。

最小化 CUDA warp 分歧效应

在单指令，多线程（SIMT）执行模型中，线程被分组成 32 个线程的集合，每个集合称为warp。如果一个 warp 遇到条件语句或分支，其线程可以分歧并串行执行每个条件。这称为分支分歧，它会显著影响性能。

CUDA warp 分歧是指在 warp 中 CUDA 线程的分歧操作。如果条件分支具有if-else结构，并且 warp 具有此 warp 分歧，所有 CUDA 线程对于分支代码块都有活动和非活动操作部分。

下图显示了 CUDA warp 中的 warp 分歧效应。不处于空闲状态的 CUDA 线程会降低 GPU 线程的有效使用：

随着分支部分的增加，GPU 调度吞吐量变得低效。因此，我们需要避免或最小化这种 warp 分歧效应。您可以选择几种选项：

通过处理不同的 warp 来避免分歧效应
通过合并分支部分来减少 warp 中的分支
缩短分支部分；只有关键部分进行分支
重新排列数据（即转置，合并等）
使用协作组中的tiled_partition来对组进行分区

确定分歧作为性能瓶颈

从先前的减少优化中，您可能会发现由于计算分析中的分歧分支而导致内核效率低下的警告，如下所示：

73.4％的分歧意味着我们有一个低效的操作路径。我们可以确定减少寻址是问题所在，如下所示：

__global__ void reduction_kernel(float* d_out, float* d_in, unsigned int size) {
    unsigned int idx_x = blockIdx.x * blockDim.x + threadIdx.x;

    extern __shared__ float s_data[];
    s_data[threadIdx.x] = (idx_x < size) ? d_in[idx_x] : 0.f;

    __syncthreads();

    // do reduction
    for (unsigned int stride = 1; stride < blockDim.x; stride *= 2) {
        // thread synchronous reduction
        if ( (idx_x % (stride * 2 - 1)) == 0 )
            s_data[threadIdx.x] += s_data[threadIdx.x + stride];

        __syncthreads();
    }

    if (threadIdx.x == 0)
        d_out[blockIdx.x] = s_data[0];
}

在减少寻址方面，我们可以选择以下 CUDA 线程索引策略之一：

交错寻址
顺序寻址

让我们回顾一下它们，并通过实施这些策略来比较它们的性能。由于我们只会修改减少内核，因此我们可以重用主机代码进行下两个实现。

交错寻址

在这种策略中，连续的 CUDA 线程使用交错寻址策略获取输入数据。与之前的版本相比，CUDA 线程通过增加步幅值来访问输入数据。以下图表显示了 CUDA 线程如何与减少项交错：

可以实现以下交错寻址：

__global__ void
 interleaved_reduction_kernel(float* g_out, float* g_in, unsigned int size) {
    unsigned int idx_x = blockIdx.x * blockDim.x + threadIdx.x;

    extern __shared__ float s_data[];
    s_data[threadIdx.x] = (idx_x < size) ? g_in[idx_x] : 0.f;
    __syncthreads();

    // do reduction
    // interleaved addressing
    for (unsigned int stride = 1; stride < blockDim.x; stride *= 2) {
        int index = 2 * stride * threadIdx.x;
        if (index < blockDim.x)
            s_data[index] += s_data[index + stride];
        __syncthreads();
    }
    if (threadIdx.x == 0)
        g_out[blockIdx.x] = s_data[0];
}

运行以下命令来编译上述代码：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o reduction ./reduction.cpp ./reduction_kernel_interleaving.cu

在 Tesla V100 上，测得的内核执行时间为 0.446 毫秒。这比之前的版本慢，因为在这种方法中每个线程块都没有完全利用。通过对其指标进行分析，我们可以得到更多细节。

现在我们将尝试另一种寻址方法，该方法旨在使每个线程块计算更多数据。

顺序寻址

与之前的版本相比，这具有高度合并的索引和寻址。这种设计更有效，因为当步幅大小大于 warp 大小时就没有分歧。以下图表显示了合并的线程操作：

现在，让我们编写一个内核函数，以在减少项上使用顺序寻址。

__global__ void
 sequantial_reduction_kernel(float *g_out, float *g_in, 
                             unsigned int size)
{
    unsigned int idx_x = blockIdx.x * blockDim.x + threadIdx.x;

    extern __shared__ float s_data[];

    s_data[threadIdx.x] = (idx_x < size) ? g_in[idx_x] : 0.f;

    __syncthreads();

    // do reduction
    // sequential addressing
    for (unsigned int stride = blockDim.x / 2; stride > 0; 
         stride >>= 1)
    {
        if (threadIdx.x < stride)
            s_data[threadIdx.x] += s_data[threadIdx.x + stride];

        __syncthreads();
    }

    if (threadIdx.x == 0)
        g_out[blockIdx.x] = s_data[0];
}

运行以下命令来编译上述代码：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o reduction ./reduction.cpp ./reduction_kernel_sequential.cu

在 Tesla V100 GPU 上，测得的执行时间为 0.378 毫秒，略快于之前的策略（0.399 毫秒）。

由于避免 warp 分歧，我们可以在原始计算上获得 12.2 倍的性能提升。以下图表显示了更新后的性能限制器分析：

与之前的性能限制器相比，我们可以看到减少了控制流操作并增加了内存利用率。

性能建模和平衡限制器

根据性能限制器分析，我们当前的减少性能受到计算延迟的限制，这是由于内存带宽，尽管限制器分析显示每个资源的充分利用。让我们讨论为什么这是一个问题，以及如何通过遵循 Roofline 性能模型来解决这个问题。

Roofline 模型

Roofline 模型是一种直观的视觉性能分析模型，用于为并行处理单元上的给定计算内核提供估计性能。根据这个模型，并行编程中的开发人员可以确定算法应该受到什么限制，并确定哪些应该进行优化。

以下图表显示了 Roofline 模型的一个示例：

倾斜部分表示内存受限，平坦部分表示算术受限。每个并行算法和实现都有自己的 Roofline 模型，因为它们具有不同的计算能力和内存带宽。有了这个模型，算法可以根据它们的操作密度（flops/bytes）进行放置。如果一个实现不符合这个模型的预期性能，我们可以确定这个版本受到延迟的限制。

考虑到我们并行减少的复杂性，它必须是内存受限的。换句话说，它的操作密度低，因此我们的策略应尽可能最大化内存带宽。

因此，我们需要确认我们的减少内核函数如何使用性能分析器中的内存带宽。以下图表显示了全局内存的带宽使用情况：

如图所示，我们没有充分利用内存带宽。Tesla V100 GPU 的总带宽为 343.376 GB/s，利用了大约三分之一的带宽，因为这款 GPU 具有 900 GB/s 带宽的 HBM2 内存。因此，下一步是通过让每个 CUDA 线程处理更多数据来增加带宽使用率。这将解决延迟限制的情况，并使我们的应用程序受限于内存带宽。

现在，让我们讨论如何增加内存带宽。

通过网格跨步循环最大化内存带宽

我们可以通过一个简单的想法实现这一点。减少问题允许我们使用 CUDA 线程累积输入数据并开始减少操作。以前，我们的减少实现是从输入数据大小开始的。但现在，我们将迭代到一组 CUDA 线程的输入数据，并且该大小将是我们内核函数的网格大小。这种迭代风格称为网格跨步循环。这种技术有许多好处，可以控制多个 CUDA 核心，并在本文中介绍：devblogs.nvidia.com/cuda-pro-tip-write-flexible-kernels-grid-stride-loops。

以下代码显示了更新后的减少内核函数：

__global__ void reduction_kernel(float *g_out, float *g_in, 
                                 unsigned int size) {
    unsigned int idx_x = blockIdx.x * blockDim.x + threadIdx.x;
    extern __shared__ float s_data[];

    // cumulates input with grid-stride loop 
 // and save to the shared memory
 float input = 0.f;
 for (int i = idx_x; i < size; i += blockDim.x * gridDim.x)
 input += g_in[i];
 s_data[threadIdx.x] = input;
 __syncthreads();

    // do reduction
    for (unsigned int stride = blockDim.x / 2; stride > 0; 
         stride >>= 1) {
        if (threadIdx.x < stride)
            s_data[threadIdx.x] += s_data[threadIdx.x + stride];
        __syncthreads();
    }
    if (threadIdx.x == 0)
        g_out[blockIdx.x] = s_data[0];
}

您会发现这个内核函数首先专注于累积输入数据，然后减少加载的数据。

现在，我们需要确定网格大小。为了使我们的 GPU 代码在各种 GPU 目标上运行，我们必须在运行时确定它们的大小。此外，我们需要利用 GPU 中的所有多处理器。CUDA C 提供了相关函数。我们可以使用cudaOccpancyMaxActiveBlocksPerMultiprocessor()函数获得占用率感知的每个多处理器的最大活动块数。此外，我们可以使用cudaDeviceGetAttribte()函数获得目标 GPU 上的多处理器数量。以下代码显示了如何使用这些函数并调用内核函数：

int reduction(float *g_outPtr, float *g_inPtr, int size, int n_threads)
{
    int num_sms;
    int num_blocks_per_sm;
    cudaDeviceGetAttribute(&num_sms, 
                           cudaDevAttrMultiProcessorCount, 0);
    cudaOccupancyMaxActiveBlocksPerMultiprocessor(&num_blocks_per_sm, 
           reduction_kernel, n_threads, n_threads*sizeof(float));
    int n_blocks = min(num_blocks_per_sm * num_sms, (size 
                       + n_threads - 1) / n_threads);

    reduction_kernel<<<n_blocks, n_threads, n_threads * 
                       sizeof(float), 0>>>(g_outPtr, g_inPtr, size);
    reduction_kernel<<<1, n_threads, n_threads * sizeof(float), 
                       0>>>(g_outPtr, g_inPtr, n_blocks);
    return 1;
}

这个函数还有一个额外的修改。为了节省占用率计算开销，它再次启动reduction_kernel()函数，这次只使用一个块。运行以下命令：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o reduction ./reduction.cpp ./reduction_kernel.cu

更新后的减少性能为0.278 ms，在 Tesla V100 上比以前的方法快了大约 100 ms。

现在，让我们回顾一下我们如何利用内存带宽。以下图表显示了在 Visual Profiler 中的内存利用分析，并显示了我们如何将内存带宽增加了两倍：

尽管它显示出了增加的带宽，但我们仍然有进一步增加的空间。让我们来看看如何实现更多的带宽。

平衡 I/O 吞吐量

从分析器得到的结果来看，局部变量 input 有大量的加载/存储请求。这样大量的 I/O 会影响线程块的调度，因为存在操作依赖。当前数据累积中最糟糕的是它对设备内存有依赖。因此，我们将使用额外的寄存器来发出更多的加载指令以减轻依赖。以下代码显示了我们如何做到这一点：

#define NUM_LOAD 4
__global__ void
 reduction_kernel(float *g_out, float *g_in, unsigned int size)
{
    unsigned int idx_x = blockIdx.x * blockDim.x + threadIdx.x;

    extern __shared__ float s_data[];

    // cumulates input with grid-stride loop 
    // and save to the shared memory
    float input[NUM_LOAD] = {0.f};
    for (int i = idx_x; i < size; i += blockDim.x * 
         gridDim.x * NUM_LOAD)
    {
        for (int step = 0; step < NUM_LOAD; step++)
            input[step] += (i + step * blockDim.x * gridDim.x < size) ? 
                g_in[i + step * blockDim.x * gridDim.x] : 0.f;
    }
    for (int i = 1; i < NUM_LOAD; i++)
        input[0] += input[i];
    s_data[threadIdx.x] = input[0];

    __syncthreads();

    // do reduction
    for (unsigned int stride = blockDim.x / 2; stride > 0; 
         stride >>= 1)
    {
        if (threadIdx.x < stride)
            s_data[threadIdx.x] += s_data[threadIdx.x + stride];

        __syncthreads();
    }

    if (threadIdx.x == 0) {
        g_out[blockIdx.x] = s_data[0];
    }
}

这段代码使用了三个额外的寄存器来收集全局内存数据。NUM_LOAD的值可能会因 GPU 的不同而有所不同，因为它受 GPU 的内存带宽和 GPU 中 CUDA 核心数量的影响：

运行以下命令时，使用 Tesla V100 卡的性能达到了0.264毫秒：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o reduction ./reduction.cpp ./reduction_kernel_opt.cu

warp 级原语编程

CUDA 9.0 引入了新的 warp 同步编程。这一重大变化旨在避免 CUDA 编程依赖隐式 warp 同步操作，并明确处理同步目标。这有助于防止 warp 级同步操作中的疏忽竞争条件和死锁。

从历史上看，CUDA 只提供了一个显式同步 API，即__syncthreads()用于线程块中的 CUDA 线程，并依赖于 warp 的隐式同步。下图显示了 CUDA 线程块操作的两个级别的同步：

然而，最新的 GPU 架构（Volta 和 Turing）具有增强的线程控制模型，其中每个线程可以执行不同的指令，同时保持其 SIMT 编程模型。下图显示了它是如何改变的：

直到 Pascal 架构（左图），线程是在 warp 级别进行调度的，并且它们在 warp 内部隐式同步。因此，CUDA 线程在 warp 中隐式同步。然而，这可能会导致意外的死锁。

Volta 架构对此进行了改进，并引入了独立线程调度。这种控制模型使每个 CUDA 线程都有自己的程序计数器，并允许 warp 中的一组参与线程。在这个模型中，我们必须使用显式的同步 API 来指定每个 CUDA 线程的操作。

因此，CUDA 9 引入了显式的 warp 级原语函数：

	warp 级原语函数
识别活动线程	`__activemask()`
屏蔽活动线程	`__all_sync()`, `__any_sync()`, `__uni_sync()`, __ballot_sync()``__match_any_sync(), `__match_all_sync()`
同步数据交换	`__shfl_sync()`, `__shfl_up_sync()`, `__shfl_down_sync()`, `__shfl_xor_sync()`
线程同步	`__syncwarp()`

有三类 warp 级原语函数，分别是 warp 识别、warp 操作和同步。所有这些函数都隐式地指定了同步目标，以避免意外的竞争条件。

使用 warp 原语进行并行归约

让我们看看这如何有益于我们的并行归约实现。这个示例将使用 Cooperative Groups 中的shfl_down()函数和 warp 原语函数中的shfl_down_sync()。下图显示了shfl_down_sync()如何与 shift down 操作一起工作：

在这个集体操作中，warp 中的 CUDA 线程可以将指定的寄存器值移动到同一个 warp 中的另一个线程，并与其同步。具体来说，集体操作有两个步骤（第三个是可选的）：

识别、屏蔽或投票源 CUDA 线程在一个 warp 中将进行操作。
让 CUDA 线程移动数据。
warp 中的所有 CUDA 线程都在同步（可选）。

对于并行归约问题，我们可以使用__shfl_down_sync()进行 warp 级别的归约。现在，我们可以通过以下图来增强我们的线程块级别的归约：

每个 warp 的归约结果都存储在共享内存中，以与其他 warp 共享。然后，通过再次进行 warp-wise 收集，可以获得最终的块级归约。

我们使用__shfl_down_sync()，因为我们只需要一个线程进行 warp 级别的归约。如果您需要让所有 CUDA 线程都进行 warp 级别的归约，可以使用__shfl_xor_sync()。

第一个块级别的归约的数量是网格的维度，输出存储在全局内存中。通过再次调用，我们可以使用 warp 级别的同步函数构建一个并行归约核。

现在，让我们使用 warp 级别的原始函数来实现 warp 级别的归约。首先，我们将编写一个使用 warp-shifting 函数进行 warp 级别归约的函数。以下代码显示了如何实现这一点：

__inline__ __device__ float warp_reduce_sum(float val) {
    for (int offset = warpSize / 2; offset > 0; offset >>= 1) {
        unsigned int mask = __activemask();
        val += __shfl_down_sync(mask, val, offset);
    }
    return val;
}

对于 warp-shifting，我们需要让 CUDA 调度程序识别活动线程，并让 warp-shifting 函数进行归约。

第二步是使用先前的 warp 级别归约编写一个块级别的归约函数。我们将在共享内存中收集先前的结果，并从结果中进行第二次归约。以下代码显示了如何实现这一点：

__inline__ __device__ float block_reduce_sum(float val) {
    // Shared mem for 32 partial sums
    static __shared__ float shared[32]; 
    int lane = threadIdx.x % warpSize;
    int wid = threadIdx.x / warpSize;

    val = warp_reduce_sum(val); // Warp-level partial reduction
    if (lane == 0)
        shared[wid] = val; // Write reduced value to shared memory
    __syncthreads(); // Wait for all partial reductions

    //read from shared memory only if that warp existed
    if (wid == 0) {
        val = (threadIdx.x < blockDim.x / warpSize) ? shared[lane] : 0;
        val = warp_reduce_sum(val); //Final reduce within first warp
    }
    return val;
}

现在，我们将实现归约核函数，累积输入数据，并从我们实现的块级归约中进行归约。由于我们只关注优化 warp 级别的优化，因此整体设计与之前的版本相同。以下代码显示了核函数：

__global__ void
reduction_kernel(float *g_out, float *g_in, unsigned int size) {
    unsigned int idx_x = blockIdx.x * blockDim.x + threadIdx.x;
    // cumulates input with grid-stride loop and save to share memory
    float sum[NUM_LOAD] = { 0.f };
    for (int i = idx_x; i < size; i += blockDim.x * gridDim.x * NUM_LOAD) {
        for (int step = 0; step < NUM_LOAD; step++)
            sum[step] += (i + step * blockDim.x * gridDim.x < size) ? g_in[i + step * blockDim.x * gridDim.x] : 0.f;
    }
    for (int i = 1; i < NUM_LOAD; i++)
        sum[0] += sum[i];
    // warp synchronous reduction
    sum[0] = block_reduce_sum(sum[0]);

    if (threadIdx.x == 0)
        g_out[blockIdx.x] = sum[0];
}

然后，让我们使用以下命令编译代码：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o reduction ./reduction.cpp ./reduction_wp_kernel.cu

以下屏幕截图显示了执行时间的减少：

在主机代码修改上，没有从 warp 原语切换到协作组。因此，我们可以对两种归约实现使用相同的主机代码。

我们已经介绍了 CUDA 中的 warp 同步编程。它的应用不仅限于归约，还可以用于其他并行算法：扫描、双调排序和转置。如果您需要了解更多信息，可以查看以下文章：

灵活处理线程的协作组

CUDA 9.0 引入了一个名为协作组的新 CUDA 编程特性。这通过指定组操作来引入了一种新的 CUDA 编程设计模式。使用这个特性，程序员可以编写显式控制 CUDA 线程的 CUDA 代码。

首先，让我们看看协作组是什么，以及它的编程优势。

CUDA 线程块中的协作组

协作组提供了显式的 CUDA 线程分组对象，帮助程序员更清晰、更方便地编写集体操作。例如，我们需要获取一个掩码来控制 warp 中活动的 CUDA 线程以进行 warp-shifting 操作。另一方面，协作组对象将可用的线程绑定为一个瓦片，并将它们作为一个对象进行控制。这为 CUDA C 编程带来了 C++语言的好处。

协作组的基本类型是thread_group。这使得 C++类样式的类型thread_group能够提供其配置信息，使用is_valid()、size()和thread_rank()函数。此外，这提供了可以应用于组中所有 CUDA 线程的集体函数。这些函数如下：

	thread_group 集体函数
识别活动线程	`tiled_partition()`, `coalesced_threads()`
屏蔽活动线程	`any()`, `all()`, ballot()``match_any(), `match_all()`
同步数据交换	`shfl()`, `shfl_up()`, `shfl_down()`, `shfl_xor()`
线程同步	`sync()`

这些函数列表类似于 warp 级别的原始函数。因此，warp 级别的原始操作可以用协作组替换。thread_group可以被较小的thread_group、thread_block_tile或coalesced_group分割。

协作组还提供了线程块编程的灵活性。使用以下代码行，我们可以处理一个线程块：

thread_block block = this_thread_block();

thread_block提供了 CUDA 内置关键字包装函数，我们使用它来获取块索引和线程索引：

dim3 group_index();  // 3-dimensional block index within the grid
dim3 thread_index(); // 3-dimensional thread index within the block

我们可以使用this_thread_block()来获取一个线程块对象，如下所示：

thread_block block = this_thread_block();

现在，让我们看看协作组的好处与传统的 CUDA 内置变量相比有什么好处。

协作组的好处

使用协作组提供了更多的 C++可编程性，而不是使用传统的 CUDA 内置变量。使用thread_block组，您可以将您的内核代码从使用内置变量切换到协作组的索引。但是，协作组的真正力量不仅仅如此。让我们在以下部分介绍其优势。

模块化

使用协作组，程序员可以将集体操作的内核代码模块化，对应于屏障目标。这有助于避免假设所有线程都在同时运行而导致的疏忽，从而引发死锁和竞争条件。以下是 CUDA 线程同步的死锁和正常操作的示例：

对于左侧的示例，内核代码意图同步 CUDA 线程块中的一部分线程。通过指定屏障目标，此代码最小化了同步开销。然而，它引入了死锁情况，因为__syncthreads()调用了一个屏障，等待所有 CUDA 线程到达屏障。然而，__synchthroead()无法满足其他线程的要求并等待。右侧的示例显示了良好的操作，因为它没有任何死锁点，因为线程块中的所有线程都可以满足__syncthreads()。

在协作组 API 中，CUDA 程序员指定线程组进行同步。协作组使得显式同步目标成为可能，因此程序员可以让 CUDA 线程显式同步。这个项目也可以被视为一个实例，因此我们可以将实例传递给设备函数。

以下代码显示了协作组如何提供显式同步对象并将它们作为实例处理：

__device__ bar(thread_group block, float *x) {
    ...
    block.sync();
    ...
}

__global__ foo() {
    bar(this_thread_block(), float *x);
}

正如在前面的示例代码中所示，内核代码可以指定同步组并将它们作为thread_group参数传递。这有助于我们在子例程中指定同步目标。因此，程序员可以通过使用协作组来防止意外死锁。此外，我们可以将不同类型的组设置为thread_group类型并重用同步代码。

显式分组线程的操作和避免竞争条件

协作组通过在 warp 中平铺线程来支持 warp 级协作操作。如果 tile 大小与 warp 大小匹配，CUDA 可以省略 warps 的隐式同步，确保正确的内存操作以避免竞争条件。通过消除隐式同步，可以增强 GPU 的性能。从历史上看，有经验的 CUDA 程序员使用分离的 warps 进行 warp 级同步。这意味着 warp 中的协作操作不必与其他 warp 操作同步。这释放了 GPU 的性能。但是，这是有风险的，因为它引入了协作操作之间的竞争条件。

动态活动线程选择

CUDA Cooperative Groups 的另一个好处是程序员可以选择 warp 中的活动线程，以避免分支分歧效应。由于 CUDA 是 SIMT 架构，一个指令单元发出一组线程，并且如果它们遇到分支，就无法禁用分歧。但是，从 CUDA 9.0 开始，程序员可以使用coalesced_threads()选择在分支块中活动的线程。这通过禁用不参与分支的线程返回聚合的线程。然后，SM 的指令单元发出下一个活动线程组中的活动线程。

应用于并行减少

我们将更新以前的减少内核代码以使用协作组。从以前的内核代码中，您可以轻松应用协作组的thread_block，如下所示：

__global__ void
 reduction_kernel(float* g_out, float* g_in, unsigned int size)
{
    unsigned int idx_x = blockIdx.x * blockDim.x + threadIdx.x;

    thread_block block = this_thread_block();

    extern __shared__ float s_data[];

我们不必更新数据输入累积部分，因此让我们为每个线程块更新减少部分。以下代码显示了一个块大小的减少的示例：

    // do reduction
    for (unsigned int stride = block.group_dim().x / 2; stride > 0; 
         stride >>= 1) {
        if (block.thread_index().x < stride) {
            s_data[block.thread_index().x] += 
                s_data[block.thread_index().x + stride];
            block.sync(); // threads synchronization in a branch
        }
    }
}

使用以下命令的估计操作性能为 0.264 毫秒：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o reduction_cg -rdc=true ./reduction.cpp ./reduction_cg_kernel.cu

前面的命令显示了与以前版本相同的性能。

协作组以避免死锁

协作组可以支持独立的 CUDA 线程调度。因此，我们可以使用一个组单独控制 CUDA 线程，并显式地对它们进行同步。目标组可以是预定义的 tile，但也可以根据条件分支确定，如下所示：

// do reduction
for (unsigned int stride = block.group_dim().x / 2; stride > 0; 
     stride >>= 1) {
    // scheduled threads reduce for every iteration
    // and will be smaller than a warp size (32) eventually.
    if (block.thread_index().x < stride) { 
        s_data[block.thread_index().x] += s_data[
                       block.thread_index().x + stride];

        // __syncthreads(); // (3) Error. Deadlock.
        // block.sync();    // (4) Okay. Benefit of Cooperative Group
    }
    // __syncthreads();     // (1) Okay
    block.sync();           // (2) Okay
}

这段代码有四种线程块同步选项。选项(1)和(2)是具有不同 API 的等效操作。另一方面，选项(3)和(4)则不是。选项(3)引入了 CUDA 线程的死锁，主机无法返回 CUDA 内核，因为活动的 CUDA 线程无法与未激活的 CUDA 线程同步。另一方面，选项(4)由于协作组的自动活动线程识别而起作用。这有助于我们避免意外错误并轻松开发复杂的算法。

NVIDIA 提供了有关协作组的详细描述，可以在以下文档中找到：

您还可以从cooperative_groups.h本身了解其架构和完整的 API 列表。

CUDA 内核中的循环展开

CUDA 也可以像其他编程语言一样受益于循环展开。通过这种技术，CUDA 线程可以减少或消除循环控制开销，例如循环结束测试每次迭代，分支惩罚等。

如果 CUDA 可以识别循环的迭代次数，它会自动展开小循环。程序员还可以使用#pragma unroll指令向编译器提供提示，或者将循环代码重写为一组独立的语句。应用循环展开很简单，因此您可以轻松应用到当前的工作代码中。

让我们将这应用到我们的并行减少实现中。就像 C/C++中的普通循环展开指令一样，我们可以在for循环的顶部放置#pragma循环展开指令。NVCC 编译器可以展开循环，因为编译器可以自行获得group.size()的确切大小：

template <typename group_t>
__inline__ __device__ float
 warp_reduce_sum(group_t group, float val)
{
    #pragma unroll
    for (int offset = group.size() / 2; offset > 0; offset >>= 1)
        val += group.shfl_down(val, offset);
    return val;
}

使用以下命令，估计的操作性能为 0.263 毫秒：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o reduction_cg -rdc=true ./reduction.cpp ./reduction_cg_kernel.cu

如果您更喜欢使用 warp 原始函数，可以像下面这样编写warp_reduce_sum。循环代码可以通过用warpSize替换group.size()来重用，但在这种情况下稍微更快：

#define FULL_MASK 0xFFFFFFFF
__inline__ __device__ float
warp_reduce_sum(float val) {
#pragma unroll 5
    for (int offset = 1; offset < 6; offset++)
        val += __shfl_down_sync(FULL_MASK, val, warpSize >> offset);
    return val;
}

运行以下命令来编译上述代码：

nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o reduction_wp -rdc=true ./reduction.cpp ./reduction_wp_kernel.cu

其结果是 0.263 毫秒，与先前的结果相同。

使用循环展开存在一个陷阱。展开的代码执行可能导致寄存器使用增加而降低占用率。此外，由于代码执行大小增加，可能会出现更高的指令缓存未命中惩罚。

原子操作

在 CUDA 编程中，程序员可以使用原子 API 从多个 CUDA 线程更新共享资源。这些原子 API 保证消除对共享资源的竞争条件，因此我们可以期望并行执行产生一致的输出。这个操作对于获取统计参数（如直方图、均值、总和等）特别有用。我们还可以简化代码实现。例如，可以使用以下代码中的atomicAdd()函数编写减少操作：

__global__ void
 atomic_reduction_kernel(float *data_out, float *data_in, int size)
 {
     int idx_x = blockIdx.x * blockDim.x + threadIdx.x;
     atomicAdd(&data_out[0], data_in[idx_x]);
 }

正如您所看到的，原子函数简化了所需的操作。但是，由于原子操作将所有请求串行化到共享资源，因此其性能较慢。运行以下命令查看执行时间：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o mixed_precision_single ./mixed_precision.cu

这个显示的内核函数在我的 Tesla V100 上花费了 39 毫秒，比原始版本（4.609 毫秒）慢得多。因此，建议的原子操作使用是只在必要时限制请求。例如，对于并行减少问题，我们可以在某个级别并行减少项目，并使用原子操作输出最终结果。

下图显示了另一种可能的方法。这将块内减少替换为atomicAdd：

在上图中，我们可以看到有两个减少点：一个是warp，一个是线程块，并且，块内减少结果通过单个全局内存变量原子地累积。因此，我们可以消除第二次减少迭代。以下截图显示了第二次减少迭代的内核优化优先级（左侧）和性能限制分析（右侧）：

内核优化优先级与性能限制分析（第二次迭代）

换句话说，第二次迭代的性能受其较小的网格大小的延迟限制。因此，通过删除这一点，我们将能够减少执行时间。

现在让我们实现该设计并看看性能如何改变。我们只需要更新减少内核函数的最后部分：

__global__ void
 reduction_kernel(float* g_out, float* g_in, unsigned int size)
{
    unsigned int idx_x = blockIdx.x * (2 * blockDim.x) + threadIdx.x;

    thread_block block = this_thread_block();

    // cumulates input with grid-stride loop and save to share memory
    float sum[NUM_LOAD] = { 0.f };
    for (int i = idx_x; i < size; i += blockDim.x 
         * gridDim.x * NUM_LOAD)
    {
        for (int step = 0; step < NUM_LOAD; step++)
            sum[step] += (i + step * blockDim.x * gridDim.x < size) ? 
                         g_in[i + step * blockDim.x * gridDim.x] : 0.f;
    }
    for (int i = 1; i < NUM_LOAD; i++)
        sum[0] += sum[i];
    // warp synchronous reduction
    sum[0] = block_reduce_sum(block, sum[0]);

    sum[0] = block_reduce_sum(sum[0]);

    // Performing Atomic Add per block
    if (block.thread_rank() == 0) {
        atomicAdd(&g_out[0], sum);
    }
}

然后，我们将删除第二次迭代函数调用。因此，如果原子操作的延迟短于那个，我们可以消除内核调用延迟并实现更好的性能。运行以下命令：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o reduction_atomic_block ./reduction.cpp ./reduction_blk_atmc_kernel.cu

幸运的是，在 Tesla V100 上估计的执行时间为 0.259 毫秒，因此我们可以获得稍微增强的结果。

如果您想了解 CUDA C 中原子操作的更多信息，请查看此链接的编程指南：docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#atomic-functions。

低/混合精度操作

混合精度是一种探索低精度并获得高精度结果的技术。这种技术使用低精度计算核心操作，并使用高精度操作生成输出。与高精度计算相比，低精度操作计算具有减少内存带宽和更高的计算吞吐量的优势。如果低精度足以从具有高精度的应用程序中获得目标精度，这种技术可以通过这种权衡来提高性能。NVIDIA 开发者博客介绍了这种可编程性：devblogs.nvidia.com/mixed-precision-programming-cuda-8。

在这些情况下，CUDA 将其支持扩展到低于 32 位数据类型的低精度工具，例如 8/16 位整数（INT8/INT16）和 16 位浮点数（FP16）。对于这些低精度数据类型，GPU 可以使用一些特定的 API 进行单指令，多数据（SIMD）操作。在本节中，我们将研究这两种用于混合精度目的的低精度操作的指令。

要从中受益，您需要确认您的 GPU 是否支持低混合精度操作和支持的数据类型。特定 GPU 支持低精度计算是可能的，精度取决于 GPU 芯片组。具体来说，GP102（Tesla P40 和 Titan X），GP104（Tesla P4）和 GP106 支持 INT8；GP100（Tesla P100）和 GV100（Tesla V100）支持 FP16（半精度）操作。Tesla GV100 兼容 INT8 操作，没有性能下降。

CUDA 具有一些特殊的内置函数，可以为低精度数据类型启用 SIMD 操作。

半精度操作

CUDA 为半精度浮点数据类型（FP16）提供了内置函数，并且开发人员可以选择 CUDA 是否为每条指令计算一个或两个值。CUDA 还提供了单精度和半精度之间的类型转换函数。由于 FP16 的精度限制，您必须使用转换内置函数来处理单精度值。

现在，让我们实现和测试 GPU 的 FP16 操作。GPU 可以支持高于计算能力 5.3 的这种类型的本机计算。但是有些 GPU 不支持这一点，因此请仔细检查您的 GPU 是否支持这种半精度操作。

CUDA C 中的半精度数据类型是half，但您也可以使用__half类型。对于 API，CUDA 提供了与此数据类型相关的内置函数，例如__hfma()、__hmul()和__hadd()。这些内置函数还提供了使用__hfma2()、__hmul2()和__hadd2()一次处理两个数据的本机操作。使用这些函数，我们可以编写混合精度操作的核心代码：

__global__ void hfma_kernel(half *d_x, half *d_y, float *d_z, int size)
 {
     int idx_x = blockIdx.x * blockDim.x + threadIdx.x;
     int stride = gridDim.x * blockDim.x;

     half2 *dual_x = reinterpret_cast<half2*>(d_x);
     half2 *dual_y = reinterpret_cast<half2*>(d_y);
     float2 *dual_z = reinterpret_cast<float2*>(d_z);

     extern __shared__ float2 s_data[];

 #if __CUDA_ARCH__ >= 530
     for (int i = idx_x; i < size; i+=stride) {
         s_data[threadIdx.x] = __half22float2(__hmul2(dual_y[i], 
                                                      dual_x[i]));
         __syncthreads();
         dual_z[i] = s_data[threadIdx.x];
     }
     #else
     for (int i = idx_x; i < size; i+=stride) {
         s_data[threadIdx.x] = __half22float2(dual_x[i]) * 
                               __half22float2(dual_y[i]);
         __syncthreads();
         dual_z[i] = s_data[threadIdx.x];
     }
     #endif
 }

对于那些不支持本机半精度操作的 GPU，我们的代码在编译时检查 CUDA 的计算能力，并确定应采取哪种操作。

以下代码调用了半精度网格大小的核函数，因为每个 CUDA 线程将操作两个数据：

int n_threads = 256;
int num_sms;
int num_blocks_per_sm;
cudaDeviceGetAttribute(&num_sms, cudaDevAttrMultiProcessorCount, 0);
cudaOccupancyMaxActiveBlocksPerMultiprocessor(&num_blocks_per_sm,   
    hfma_kernel, n_threads, n_threads*sizeof(float2));
int n_blocks = min(num_blocks_per_sm * num_sms, 
                   (size/2 + n_threads - 1) / n_threads);
hfma_kernel<<< n_blocks, n_threads, n_threads * sizeof(float2) 
               >>>(X.d_ptr_, Y.d_ptr_, Z.d_ptr_, size/2);

其他初始化代码和基准代码在示例配方代码中实现，因此请查看它。

我们已经涵盖了 FP16 精度操作中的 FMA 操作。CUDA C 提供了各种半精度操作（docs.nvidia.com/cuda/cuda-math-api/group__CUDA__MATH__INTRINSIC__HALF.html）。请查看其他操作。

8 位整数和 16 位数据的点积运算和累加（DP4A 和 DP2A）

对于 8 位/16 位整数，CUDA 提供了矢量化的点积操作。这些是 DP4A（四元素点积累加）和 DP2A（两元素点积累加）。使用这些函数，CUDA 开发人员可以进行更快的操作。CUDA 8.0 开发博客通过直观的图示介绍了这些函数 (devblogs.nvidia.com/mixed-precision-programming-cuda-8/)。以下显示了 GPU 的点积和累加操作的工作原理：

使用这个，你可以编写只有 8 位或 8 位/16 位混合操作的 32 位整数累加。其他操作，如求和、加法和比较，也可以使用 SIMD 内在函数。

如前所述，有特定的 GPU 可以支持 INT8/INT16 操作，具有特殊功能（dp4a和dp2a）。支持的 GPU 的计算能力必须高于 6.1。

现在，让我们实现一个使用dp4aAPI 的内核函数，如下所示：

__global__ void dp4a_kernel(char *d_x, char *d_y, int *d_z, int size)
 {
     int idx_x = blockIdx.x * blockDim.x + threadIdx.x;
     int stride = gridDim.x * blockDim.x;

 #if __CUDA_ARCH__ >= 610
     char4 *quad_x = (char4 *)d_x;
     char4 *quad_y = (char4 *)d_y;

     for (int i = idx_x; i < size; i+=stride)
         d_z[i] = __dp4a(quad_y[i], quad_x[i], 0);
 #else
     for (int i = idx_x; i < size; i+=4*stride) {
         int sum = 0;
         for (int j = 0; j < 4; j++)
             sum += d_y[4 * i + j] * d_x[4 * i + j];
         d_z[i] = sum + 0;
     }
 #endif
 }

在这个函数中，__dp4a获取两个字符数组，合并四个项目，并输出其点积输出。自帕斯卡以来，这个 API 就得到了支持，具有 CUDA 计算能力（版本 6.1）。但是旧的 GPU 架构，低于版本 6.1，需要使用原始操作。

以下代码显示了我们将如何调用实现的内核函数。由于每个 CUDA 线程将操作四个项目，其网格大小减小了四倍：

int n_threads = 256;
int num_sms;
int num_blocks_per_sm;
cudaDeviceGetAttribute(&num_sms, cudaDevAttrMultiProcessorCount, 0);
cudaOccupancyMaxActiveBlocksPerMultiprocessor(&num_blocks_per_sm, 
    dp4a_kernel, n_threads, n_threads*sizeof(int));
int n_blocks = min(num_blocks_per_sm * num_sms, (size/4 + n_threads 
                                                  - 1) / n_threads);
dp4a_kernel<<< n_blocks, n_threads, n_threads * sizeof(int) >>>  
    (X.d_ptr_, Y.d_ptr_, Z.d_ptr_, size/4);

其他初始化代码和基准代码都在示例代码中实现，就像前面的示例代码一样。

我们已经介绍了 INT8 的点操作，但 CUDA C 还提供了其他 INT8 类型的 SIMD 内在函数（docs.nvidia.com/cuda/cuda-math-api/group__CUDA__MATH__INTRINSIC__SIMD.html）。请查阅此文档以了解其他操作。

性能测量

示例代码有三个混合精度操作的版本：单精度、半精度和 INT8。随着精度的降低，我们可以为每个 CUDA 线程添加更多的操作。

运行以下命令进行单精度、半精度和 INT8 操作：

# Single-precision
$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o mixed_precision_single ./mixed_precision.cu

# Half-precision
$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o mixed_precision_half ./mixed_precision_half.cu

# INT8 
$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o mixed_precision_int ./mixed_precision_int.cu

以下表格显示了每种精度操作的估计性能：

精度	测得的性能
FP32	59.441 GFlops
FP16	86.037 GFlops
INT8	196.225 Gops

由于我们的实现没有经过优化，所以测得的性能比 Tesla V100 的理论性能要低得多。当你对它们进行分析时，它们会报告它们的内存绑定性很高。换句话说，我们需要优化它们，使它们在算术上受限，以接近理论性能。

总结

在本章中，我们介绍了如何配置 CUDA 并行操作并对其进行优化。为了做到这一点，我们必须了解 CUDA 的分层体系结构线程块和流多处理器之间的关系。通过一些性能模型——占用率、性能限制分析和 Roofline 模型——我们可以优化更多性能。然后，我们介绍了一些新的 CUDA 线程可编程性，合作组，并学习了如何简化并行编程。我们优化了并行减少问题，并在元素中实现了 0.259 毫秒，这是与相同 GPU 相比速度提高了 17.8。最后，我们了解了 CUDA 的半精度（FP16）和 INT8 精度的 SIMD 操作。

本章我们的经验集中在 GPU 的并行处理级别编程上。然而，CUDA 编程包括系统级编程。基本上，GPU 是额外的计算资源，独立于主机工作。这增加了额外的计算能力，但另一方面也可能引入延迟。CUDA 提供了可以利用这一点并隐藏延迟并实现 GPU 的全面性能的 API 函数。我们将在下一章中介绍这一点。

第四章：核心执行模型和优化策略

CUDA 编程有一个主机操作的过程。例如，我们需要分配全局内存，将数据传输到 GPU，执行核心函数，将数据传输回主机，清理全局内存。这是因为 GPU 是系统中的一个额外处理单元，所以我们需要关心它的执行和数据传输。这是与 CPU 编程相比另一个不同的 GPU 编程方面。

在本章中，我们将涵盖 CUDA 核心执行模型和 CUDA 流，它们控制 CUDA 操作。然后，我们将讨论系统级别的优化策略。接下来，我们将涵盖 CUDA 事件来测量 GPU 事件时间，以及如何使用 CUDA 事件来测量核心执行时间。之后，我们将涵盖各种 CUDA 核心执行模型，并讨论这些特性对 GPU 操作的影响。

本章将涵盖以下主题：

使用 CUDA 流的核心执行
流水线化 GPU 执行
CUDA 回调函数
具有优先级的 CUDA 流
使用 CUDA 事件估计核心执行时间
CUDA 动态并行性
网格级协作组
使用 OpenMP 的 CUDA 核心调用
多进程服务
核心执行开销比较

技术要求

本章要求我们使用的 CUDA 版本应该晚于 9.x，并且 GPU 架构应该是 Volta 或 Turing。如果你使用的是 Pascal 架构的 GPU，那么跳过Grid-level cooperative groups部分，因为这个特性是为 Volta 架构引入的。

使用 CUDA 流的核心执行

在 CUDA 编程中，流是与 GPU 相关的一系列命令。换句话说，所有的核心调用和数据传输都由 CUDA 流处理。默认情况下，CUDA 提供了一个默认流，所有的命令都隐式地使用这个流。因此，我们不需要自己处理这个。

CUDA 支持显式创建额外的流。虽然流中的操作是顺序的，但 CUDA 可以通过使用多个流同时执行多个操作。让我们学习如何处理流，以及它们具有哪些特性。

CUDA 流的使用

以下代码展示了如何创建、使用和终止 CUDA 流的示例：

cudaStream_t stream;
cudaStreamCreate(&stream);
foo_kernel<<< grid_size, block_size, 0, stream >>>();
cudaStreamDestroy(stream);

正如你所看到的，我们可以使用cudaStream_t来处理 CUDA 流。而且，我们可以使用cudaStreamCreate()来创建它，使用cudaStreamDestroy()来终止它。注意我们应该提供一个指向cudaStreamCreate()的指针。创建的流会传递给核心函数的第四个参数。

然而，我们之前并没有提供这样的流。这是因为 CUDA 提供了一个默认流，以便所有的 CUDA 操作都可以进行。现在，让我们编写一个使用默认流和多个流的应用程序。然后，我们将看到我们的应用程序如何改变。

首先，让我们编写一个使用默认 CUDA 流的应用程序，如下所示：

__global__ void foo_kernel(int step)
{
    printf("loop: %d\n", step);
}

int main()
{
    for (int i = 0; i < 5; i++)
 // CUDA kernel call with the default stream
 foo_kernel<<< 1, 1, 0, 0 >>>(i);
    cudaDeviceSynchronize();
    return 0;
}

正如你在代码中看到的，我们以流 ID 为0调用了核心函数，因为默认流的标识值为0。编译代码并查看执行输出：

$ nvcc -m64 -run -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o cuda_default_stream ./1_cuda_default_stream.cu

输出是什么？我们可以期待输出将是循环索引的顺序。以下时间轴视图显示了这段代码的操作：

可以预期，在同一个流中进行循环操作将显示核心执行的顺序。那么，如果我们使用多个 CUDA 流，并且每个循环步骤使用不同的流，会有什么改变？以下代码展示了使用不同流从 CUDA 核心函数打印循环索引的示例：

__global__ void foo_kernel(int step)
{
    printf("loop: %d\n", step);
}

int main()
{
    int n_stream = 5;
    cudaStream_t *ls_stream;
    ls_stream = (cudaStream_t*) new cudaStream_t[n_stream];

    // create multiple streams
    for (int i = 0; i < n_stream; i++)
        cudaStreamCreate(&ls_stream[i]);

    // execute kernels with the CUDA stream each
    for (int i = 0; i < n_stream; i++)
        foo_kernel<<< 1, 1, 0, ls_stream[i] >>>(i);

    // synchronize the host and GPU
    cudaDeviceSynchronize();

    // terminates all the created CUDA streams
    for (int i = 0; i < n_stream; i++)
        cudaStreamDestroy(ls_stream[i]);
    delete [] ls_stream;

    return 0;
}

在这段代码中，我们有五个调用，与之前的代码相同，但这里我们将使用五个不同的流。为此，我们建立了一个cudaStream_t数组，并为每个流创建了流。你对这个改变有什么期待？打印输出将与之前的版本相同。运行以下命令来编译这段代码：

$ nvcc -m64 -run -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o cuda_mutli_stream ./2_cuda_multi_stream.cu

然而，这并不保证它们具有相同的操作。正如我们在开始时讨论的，这段代码展示了多个流的并发性，如下面的截图所示：

正如你在截图底部所看到的，五个独立的流同时执行相同的内核函数，并且它们的操作相互重叠。由此，我们可以得出流的两个特点，如下所示：

内核执行与主机是异步的。
不同流中的 CUDA 操作是彼此独立的。

利用流的并发性，我们可以通过重叠独立操作来获得额外的优化机会。

流级别的同步

CUDA 流提供了流级别的同步，使用cudaStreamSynchronize()函数。使用这个函数会强制主机等待直到某个流的操作结束。这为我们迄今为止使用的cudaDeviceSynchronize()函数提供了重要的优化。

我们将在接下来的部分讨论如何利用这一特性，但让我们在这里讨论它的基本操作。前面的例子展示了在循环中没有同步的并发操作。然而，我们可以通过使用cudaStreamSynchronize()函数来阻止主机执行下一个内核执行。下面的代码展示了在内核执行结束时使用流同步的示例：

// execute kernels with the CUDA stream each
for (int i = 0; i < n_stream; i++) {
   foo_kernel<<< 1, 1, 0, ls_stream[i] >>>(i);
   cudaStreamSynchronize(ls_stream[i]);
}

我们可以很容易地预测，由于同步，内核操作的并发性将消失。为了确认这一点，让我们对此进行分析，看看这对内核执行的影响：

$ nvcc -m64 -run -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o cuda_mutli_stream_with_sync ./3_cuda_multi_stream_with_sync.cu

下面的截图显示了结果：

正如你所看到的，所有的内核执行没有重叠点，尽管它们是用不同的流执行的。利用这一特性，我们可以让主机等待特定流操作的开始和结果。

使用默认流

为了让多个流同时运行，我们应该使用我们显式创建的流，因为所有流操作都与默认流同步。下面的截图显示了默认流的同步操作效果：

我们可以通过修改我们的多流内核调用操作来实现这一点，就像这样：

for (int i = 0; i < n_stream; i++)
    if (i == 3)
        foo_kernel<<< 1, 1, 0, 0 >>>(i);
    else
        foo_kernel<<< 1, 1, 0, ls_stream[i] >>>(i);

运行以下命令编译代码：

$ nvcc -m64 -run -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o cuda_multi_stream_with_default ./4_cuda_multi_stream_with_default.cu

因此，我们可以看到最后一个操作无法与前面的内核执行重叠，而是必须等到第四个内核执行完成后才能进行。

GPU 执行的流水线

多个流的主要好处之一是将数据传输与内核执行重叠。通过重叠内核操作和数据传输，我们可以隐藏数据传输开销并提高整体性能。

GPU 流水线的概念

当我们执行内核函数时，我们需要将数据从主机传输到 GPU，然后将结果从 GPU 传输回主机。下面的图表显示了在主机和内核执行之间传输数据的迭代操作的示例：

然而，内核执行基本上是异步的，主机和 GPU 可以同时运行。如果主机和 GPU 之间的数据传输具有相同的特性，我们就能够重叠它们的执行，就像我们在前面的部分中看到的那样。下面的图表显示了当数据传输可以像正常的内核操作一样执行，并与流一起处理时的操作：

在这个图表中，我们可以看到主机和设备之间的数据传输可以与内核执行重叠。然后，这种重叠操作的好处是减少应用程序的执行时间。通过比较两张图片的长度，您将能够确认哪个操作的吞吐量更高。

关于 CUDA 流，所有 CUDA 操作——数据传输和内核执行——在同一个流中是顺序的。然而，它们可以与不同的流同时操作。以下图表显示了多个流的重叠数据传输和内核操作：

为了实现这样的流水线操作，CUDA 有三个先决条件：

主机内存应该分配为固定内存——CUDA 提供了cudaMallocHost()和cudaFreeHost()函数来实现这一目的。
在主机和 GPU 之间传输数据而不阻塞主机——CUDA 提供了cudaMemcpyAsync()函数来实现这一目的。
管理每个操作以及不同的 CUDA 流，以实现并发操作。

现在，让我们编写一个简单的应用程序来对工作负载进行流水线处理。

构建流水线执行

以下代码显示了异步数据传输的片段以及在执行结束时 CUDA 流的同步：

cudaStream_t stream;
float *h_ptr, *d_ptr;    size_t byte_size = sizeof(float) * BUF_SIZE;

cudaStreamCreate(&stream);               // create CUDA stream
cudaMallocHost(h_ptr, byte_size);        // allocates pinned memory
cudaMalloc((void**)&d_ptr, byte_size);   // allocates a global memory

// transfer the data from host to the device asynchronously
cudaMemcpyAsync(d_ptr, h_ptr, byte_size, cudaMemcpyHostToDevice, stream);

... { kernel execution } ...

// transfer the data from the device to host asynchronously
cudaMemcpyAsync(h_ptr, d_ptr, byte_size, cudaMemcpyDeviceToHost, stream);
cudaStreamSynchronize(stream);

// terminates allocated resources
cudaStreamDestroy(stream);
cudaFree(d_ptr);
cudaFreeHost(h_ptr);

这段代码展示了如何分配固定内存，并使用用户创建的流传输数据。通过合并这个例子和多个 CUDA 流操作，我们可以实现流水线 CUDA 操作。

现在，让我们构建一个应用程序，其中包含数据传输和内核执行的流水线操作。在这个应用程序中，我们将使用一个将两个向量相加的内核函数，通过切片流的数量，并输出其结果。然而，内核的实现在主机代码级别不需要任何更改。但是，我们将迭代加法操作 500 次以延长内核执行时间。因此，实现的内核代码如下：

__global__ void
vecAdd_kernel(float *c, const float* a, const float* b)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    for (int i = 0; i < 500; i++)
        c[idx] = a[idx] + b[idx];
}

为了处理每个流的操作，我们将创建一个管理 CUDA 流和 CUDA 操作的类。这个类将允许我们管理 CUDA 流以及索引。以下代码显示了该类的基本架构：

class Operator
{
private:
    int index;

public:
    Operator() {
        cudaStreamCreate(&stream);    // create a CUDA stream
    }

    ~Operator() {
        cudaStreamDestroy(stream);    // terminate the CUDA stream
    }

    cudaStream_t stream;
    void set_index(int idx) { index = idx; }
    void async_operation(float *h_c, const float *h_a, 
                         const float *h_b,
                         float *d_c, float *d_a, float *d_b,
                         const int size, const int bufsize);

}; // Operator

现在，让我们编写一些顺序 GPU 执行代码，这些代码在前一节中已经使用过，但作为Operator类的成员函数，如下所示：

void Operator::async_operation(float *h_c, const float *h_a, 
                          const float *h_b,
                          float *d_c, float *d_a, float *d_b,
                          const int size, const int bufsize)
{
    // start timer
    sdkStartTimer(&_p_timer);

    // copy host -> device
    cudaMemcpyAsync(d_a, h_a, bufsize, 
                    cudaMemcpyHostToDevice, stream);
    cudaMemcpyAsync(d_b, h_b, bufsize, 
                    cudaMemcpyHostToDevice, stream);

    // launch cuda kernel
    dim3 dimBlock(256);
    dim3 dimGrid(size / dimBlock.x);
    vecAdd_kernel<<< dimGrid, dimBlock, 0, 
                     stream >>>(d_c, d_a, d_b);

    // copy device -> host
    cudaMemcpyAsync(h_c, d_c, bufsize, 
                    cudaMemcpyDeviceToHost, stream);

    printf("Launched GPU task %d\n", index);
}

这个函数的操作与我们之前使用的基本 CUDA 主机编程模式没有什么不同，只是我们使用了给定的_stream应用了cudaMemcpyAsync()。然后，我们编写main()来处理多个操作符实例和页锁定内存：

int main(int argc, char* argv[])
{
    float *h_a, *h_b, *h_c;
    float *d_a, *d_b, *d_c;
    int size = 1 << 24;
    int bufsize = size * sizeof(float);
    int num_operator = 4;

    if (argc != 1)
        num_operator = atoi(argv[1]);

现在，我们将使用cudaMallocHost()来分配主机内存，以获得固定内存，并对其进行初始化：

    cudaMallocHost((void**)&h_a, bufsize);
    cudaMallocHost((void**)&h_b, bufsize);
    cudaMallocHost((void**)&h_c, bufsize);

    srand(2019);
    init_buffer(h_a, size);
    init_buffer(h_b, size);
    init_buffer(h_c, size);

而且，我们将拥有相同大小的设备内存：

    cudaMalloc((void**)&d_a, bufsize);
    cudaMalloc((void**)&d_b, bufsize);
    cudaMalloc((void**)&d_c, bufsize);

现在，我们将使用我们使用的类创建一个 CUDA 操作符列表：

    Operator *ls_operator = new Operator[num_operator];

我们准备执行流水线操作。在开始执行之前，让我们放一个秒表来查看整体执行时间，并查看重叠数据传输的好处，如下所示：

    StopWatchInterface *timer;
    sdkCreateTimer(&timer);
    sdkStartTimer(&timer);

让我们使用循环执行每个操作符，并且每个操作符将根据其顺序访问主机和设备内存。我们还将测量循环的执行时间：

    for (int i = 0; i < num_operator; i++) {
        int offset = i * size / num_operator;
        ls_operator[i].set_index(i);
        ls_operator[i].async_operation(&h_c[offset], 
                                       &h_a[offset], &h_b[offset],
                                       &d_c[offset], 
                                       &d_a[offset], &d_b[offset],
                                       size / num_operator, 
                                       bufsize / num_operator);
    }

    cudaDeviceSynchronize();
    sdkStopTimer(&timer);

最后，我们将比较一个样本的结果，并打印出整体测量性能：

    // prints out the result
    int print_idx = 256;
    printf("compared a sample result...\n");
    printf("host: %.6f, device: %.6f\n", h_a[print_idx] + 
           h_b[print_idx], h_c[print_idx]);

    // prints out the performance
    float elapsed_time_msed = sdkGetTimerValue(&timer);
    float bandwidth = 3 * bufsize * sizeof(float) / 
                      elapsed_time_msed / 1e6;
    printf("Time= %.3f msec, bandwidth= %f GB/s\n", 
           elapsed_time_msed, bandwidth);

终止句柄和内存，如下所示：

    sdkDeleteTimer(&timer);
    delete [] ls_operator;
    cudaFree(d_a);    cudaFree(d_b);    cudaFree(d_c);
    cudaFreeHost(h_a);cudaFreeHost(h_b);cudaFreeHost(h_c);

要执行代码，让我们重用前面的主机初始化函数和 GPU 内核函数。我们暂时不需要修改这些函数。使用以下命令编译代码：

$ nvcc -m64 -run -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o cuda_pipelining ./cuda_pipelining.cu

您必须使用 GPU 的计算能力版本号作为gencode选项。编译的输出如下：

Launched GPU task 0
Launched GPU task 1
Launched GPU task 2
Launched GPU task 3
compared a sample result...
host: 1.523750, device: 1.523750
Time= 29.508 msec, bandwidth= 27.291121 GB/s

正如我们所看到的，GPU 任务是按照内核执行的顺序以及流的顺序执行的。

现在，让我们来回顾一下应用程序在内部是如何运行的。默认情况下，示例代码将主机数据切片为四个部分，并同时执行四个 CUDA 流。我们可以看到每个核函数的输出以及流的执行情况。要查看重叠操作，您需要使用以下命令对执行进行分析：

$ nvprof -o overlapping_exec.nvvp ./overlapping_exec

以下截图显示了通过重叠数据传输和核函数执行来操作四个 CUDA 流：

核函数执行和数据传输之间的重叠

因此，GPU 可以忙碌直到最后一个核函数执行完成，并且我们可以隐藏大部分数据传输。这不仅增强了 GPU 的利用率，还减少了总应用程序执行时间。

在核函数执行之间，我们可以发现它们虽然属于不同的 CUDA 流，但没有争用。这是因为 GPU 调度器知道执行请求，并优先服务第一个。然而，当当前任务完成时，流多处理器可以为另一个 CUDA 流中的下一个核函数提供服务，因为它们仍然保持占用。

在多个 CUDA 流操作结束时，我们需要同步主机和 GPU，以确认 GPU 上的所有 CUDA 操作都已完成。为此，我们在循环结束后立即使用了cudaDeviceSynchronize()。此函数可以在调用点同步所选的所有 GPU 操作。

对于同步任务，我们可以用以下代码替换cudaDeviceSynchronize()函数。为此，我们还必须将私有成员_stream更改为公共成员：

for (int i = 0; i < num_operator; i++) {
    cudaStreamSynchronize(ls_operator[i]._stream);
}

当我们需要在每个流完成后从单个主机线程提供特定操作时，可以使用这个。但是，这不是一个好的操作设计，因为后续操作无法避免与其他流同步。

在循环中使用cudaStreamSynchronize()怎么样？在这种情况下，我们无法执行之前的重叠操作。以下截图显示了这种情况：

这是因为cudaStreamSynchronize()将同步每次迭代，应用程序将按顺序执行所有 CUDA 执行。在这种情况下，执行时间为 41.521 毫秒，比重叠执行时间慢了约 40%。

CUDA 回调函数

CUDA 回调函数是可调用的主机函数，由 GPU 执行上下文执行。使用此函数，程序员可以指定在 GPU 操作之后执行主机所需的主机操作。

CUDA 回调函数具有一个名为CUDART_CB的特殊数据类型，因此应该使用这种类型进行定义。使用此类型，程序员可以指定哪个 CUDA 流启动此函数，传递 GPU 错误状态，并提供用户数据。

要注册回调函数，CUDA 提供了cudaStreamAddCallback()。该函数接受 CUDA 流、CUDA 回调函数及其参数，以便从指定的 CUDA 流中调用指定的 CUDA 回调函数并获取用户数据。该函数有四个输入参数，但最后一个是保留的。因此，我们不使用该参数，它保持为0。

现在，让我们改进我们的代码，使用回调函数并输出单个流的性能。如果要分开之前的工作和这个工作，可以复制源代码。

首先，将这些函数声明放入Operator类的private区域：

StopWatchInterface *_p_timer;
static void CUDART_CB Callback(cudaStream_t stream, cudaError_t status, void* userData);
void print_time();

Callback()函数将在每个流的操作完成后被调用，print_time()函数将使用主机端计时器_p_timer报告估计的性能。函数的实现如下：

void Operator::CUDART_CB Callback(cudaStream_t stream, cudaError_t status, void* userData) {
    Operator* this_ = (Operator*) userData;
    this_->print_time();
}

void Operator::print_time() {
    sdkStopTimer(&p_timer);    // end timer
    float elapsed_time_msed = sdkGetTimerValue(&p_timer);
    printf("stream %2d - elapsed %.3f ms \n", index, 
           elapsed_time_msed);
}

为了进行正确的计时操作，我们需要在Operator类的构造函数中进行计时器初始化，并在类的终结器中进行计时器销毁。此外，我们必须在Operator::async_operation()函数的开头启动计时器。然后，在函数的末尾插入以下代码块。这允许 CUDA 流在完成先前的 CUDA 操作时调用主机端函数：

// register callback function
cudaStreamAddCallback(stream, Operator::Callback, this, 0);

现在，让我们编译并查看执行结果。您必须使用您的 GPU 的计算能力版本号作为gencode选项：

$ nvcc -m64 -run -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o cuda_callback ./cuda_callback.cu

这是我们更新的执行结果：

stream 0 - elapsed 11.136 ms
stream 1 - elapsed 16.998 ms
stream 2 - elapsed 23.283 ms
stream 3 - elapsed 29.487 ms
compared a sample result...
host: 1.523750, device: 1.523750
Time= 29.771 msec, bandwidth= 27.050028 GB/s

在这里，我们可以看到估计的执行时间以及 CUDA 流。回调函数估计其序列的执行时间。由于与其他流重叠并延迟后续 CUDA 流，我们可以看到后续 CUDA 流的执行时间延长。我们可以通过与分析结果匹配来确认这些经过的时间，如下所示：

尽管它们的测量经过时间随着流的执行而延长，但流之间的差值是固定的，我们可以从分析输出中看到这些操作。

因此，我们可以得出结论，我们可以编写主机代码，以便在每个单独的 CUDA 流操作完成后立即执行。这比从主线程同步每个流更加先进。

具有优先级的 CUDA 流

默认情况下，所有 CUDA 流具有相同的优先级，因此它们可以按正确的顺序执行其操作。此外，CUDA 流还可以具有优先级，并且可以被优先级更高的流取代。有了这个特性，我们可以有满足时间关键要求的 GPU 操作。

CUDA 中的优先级

要使用具有优先级的流，我们首先需要从 GPU 获取可用的优先级。我们可以使用cudaDeviceGetStreamPriorityRange()函数来获取这些值。它的输出是两个数值，即最低和最高的优先级值。然后，我们可以使用cudaStreamCreaetWithPriority()函数创建一个优先级流，如下所示：

cudaError_t cudaStreamCreateWithPriority(cudaStream_t* pStream, unsigned int flags, int priority)

我们应该提供两个额外的参数。第一个确定了创建的流与默认流的行为。我们可以使用cudaStreamDefault使新流与默认流同步，就像普通流一样。另一方面，我们可以使用cudaStreamNonBlocking使其与默认流并行操作。最后，我们可以在优先级范围内设置流的优先级。在 CUDA 编程中，最低值具有最高优先级。

此外，我们可以使用以下代码确认 GPU 是否支持这一点。但是，我们不必太担心这一点，因为自 CUDA 计算能力 3.5 以来，优先级流一直可用：

cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
if (prop.streamPrioritiesSupported == 0) { ... }

如果设备属性值为0，我们应该停止应用程序，因为 GPU 不支持流优先级。

具有优先级的流执行

现在，我们将重用之前带有回调的多流应用程序。在这段代码中，我们可以看到流可以按顺序操作，我们将看到如何使用优先级更改这个顺序。我们将从Operator类派生一个类，并且它将处理流的优先级。因此，我们将把成员变量流的保护级别从私有成员更改为受保护的成员。构造函数可以选择性地创建流，因为这可以由派生类完成。更改如下代码所示：

... { middle of the class Operator } ...
protected:
    cudaStream_t stream = nullptr;

public:
    Operator(bool create_stream = true) {
        if (create_stream)
            cudaStreamCreate(&stream);
        sdkCreateTimer(&p_timer);
    }
... { middle of the class Operator } ...

派生类Operator_with_priority将具有一个函数，可以根据给定的优先级手动创建一个 CUDA 流。该类的配置如下：

class Operator_with_priority: public Operator {
public:
    Operator_with_priority() : Operator(false) {}

    void set_priority(int priority) {
        cudaStreamCreateWithPriority(&stream, 
            cudaStreamNonBlocking, priority);
    }
};

当我们使用类处理每个流的操作时，我们将更新main()中的ls_operator创建代码，以使用我们之前编写的Operator_with_priority类，如下所示：

Operator_with_priority *ls_operator = new Operator_with_priority[num_operator];

当我们更新类时，这个类在我们请求它之前不会创建流。正如我们之前讨论的，我们需要使用以下代码获取 GPU 可用优先级范围：

// Get priority range
int priority_low, priority_high;
cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);
printf("Priority Range: low(%d), high(%d)\n", priority_low, priority_high);

然后，让我们创建每个操作以拥有不同的优先级流。为了简化这个任务，我们将让最后一个操作拥有最高的流，并看看 CUDA 流中的抢占是如何工作的。可以使用以下代码来实现这一点：

for (int i = 0; i < num_operator; i++) {
    ls_operator[i].set_index(i);

    // let the latest CUDA stream to have the high priority
    if (i + 1 == num_operator)
        ls_operator[i].set_priority(priority_high);
    else
        ls_operator[i].set_priority(priority_low);
}

之后，我们将执行每个操作，就像之前一样：

for (int i = 0 ; i < num_operator; i++) { 
    int offset = i * size / num_operator;
    ls_operator[i].async_operation(&h_c[offset], 
                                   &h_a[offset], &h_b[offset],
                                   &d_c[offset], 
                                   &d_a[offset], &d_b[offset],
                                   size / num_operator, 
                                   bufsize / num_operator);
}

为了获得正确的输出，让我们使用cudaDeviceSynchronize()函数同步主机和 GPU。最后，我们可以终止 CUDA 流。具有优先级的流可以使用cudaStreamDestroy()函数终止，因此在这个应用程序中我们已经做了必要的事情。

现在，让我们编译代码并查看效果。和往常一样，您需要向编译器提供正确的 GPU 计算能力版本：

$ nvcc -m64 -run -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o prioritized_cuda_stream ./prioritized_cuda_stream.cu

接下来是应用程序的输出：

Priority Range: low(0), high(-1)
stream 0 - elapsed 11.119 ms
stream 3 - elapsed 19.126 ms
stream 1 - elapsed 23.327 ms
stream 2 - elapsed 29.422 ms
compared a sample result...
host: 1.523750, device: 1.523750
Time= 29.730 msec, bandwidth= 27.087332 GB/s

从输出中，您可以看到操作顺序已经改变。Stream 3 在 Stream 1 和 Stream 2 之前。下面的屏幕截图显示了它是如何改变的：

在这个屏幕截图中，第二个 CUDA 流（在这种情况下是 Stream 19）被优先级最低的 CUDA 流（Stream 21）抢占，以便在 Stream 21 执行完毕后 Stream 19 完成其工作。请注意，数据传输的顺序不会根据这种优先级而改变。

使用 CUDA 事件估计内核执行时间

以前的 GPU 操作时间估计有一个限制，即它无法测量内核执行时间。这是因为我们在主机端使用了计时 API。因此，我们需要与主机和 GPU 同步以测量内核执行时间，考虑到对应用程序性能的开销和影响，这是不切实际的。

这可以通过使用 CUDA 事件来解决。CUDA 事件记录 GPU 端的事件以及 CUDA 流。CUDA 事件可以是基于 GPU 状态的事件，并记录调度时间。使用这个，我们可以触发以下操作或估计内核执行时间。在本节中，我们将讨论如何使用 CUDA 事件测量内核执行时间。

CUDA 事件由cudaEvent_t句柄管理。我们可以使用cudaEventCreate()创建 CUDA 事件句柄，并使用cudaEventDestroy()终止它。要记录事件时间，可以使用cudaEventRecord()。然后，CUDA 事件句柄记录 GPU 的事件时间。这个函数还接受 CUDA 流，这样我们就可以将事件时间枚举到特定的 CUDA 流。在获取内核执行的开始和结束事件之后，可以使用cudaEventElapsedTime()获取经过的时间，单位为毫秒。

现在，让我们讨论如何使用 CUDA 事件来使用这些 API。

使用 CUDA 事件

在本节中，我们将重用第二节中的多流应用程序。然后，我们使用 CUDA 事件枚举每个 GPU 内核的执行时间：

我们将使用一个简单的向量加法内核函数，如下所示：

__global__ void
vecAdd_kernel(float *c, const float* a, const float* b) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    for (int i = 0; i < 500; i++)
        c[idx] = a[idx] + b[idx];
}

这段代码有一个迭代，它延长了内核执行时间。

然后，我们将使用以下片段来测量内核执行时间。为了比较结果，我们将使用主机端的计时器和 CUDA 事件：

... { memory initializations } ...

// initialize the host timer
StopWatchInterface *timer;
sdkCreateTimer(&timer);

cudaEvent_t start, stop;
// create CUDA events
cudaEventCreate(&start);
cudaEventCreate(&stop);

// start to measure the execution time
sdkStartTimer(&timer);
cudaEventRecord(start);

// launch cuda kernel
dim3 dimBlock(256);
dim3 dimGrid(size / dimBlock.x);
vecAdd_kernel<<< dimGrid, dimBlock >>>(d_c, d_a, d_b);

// record the event right after the kernel execution finished
cudaEventRecord(stop);

// Synchronize the device to measure the execution time from the host side
cudaEventSynchronize(stop); // we also can make synchronization based on CUDA event
sdkStopTimer(&timer);

正如您在这段代码中所看到的，我们可以在内核调用之后立即记录 CUDA 事件。然而，计时器需要在 GPU 和主机之间进行同步。为了同步，我们使用cudaEventSynchronize(stop)函数，因为我们也可以使主机线程与事件同步。与此同时，这段代码只涵盖了处理计时资源和内核执行。但是，您还需要初始化所需的内存才能使其工作。

在内核执行之后，让我们编写代码报告每个计时资源的执行时间：

// print out the result
int print_idx = 256;
printf("compared a sample result...\n");
printf("host: %.6f, device: %.6f\n", h_a[print_idx] + h_b[print_idx], h_c[print_idx]);

// print estimated kernel execution time
float elapsed_time_msed = 0.f;
cudaEventElapsedTime(&elapsed_time_msed, start, stop);
printf("CUDA event estimated - elapsed %.3f ms \n", elapsed_time_msed);

现在，我们将通过终止计时资源来完成我们的应用程序，使用以下代码：

// delete timer
sdkDeleteTimer(&timer);

// terminate CUDA events
cudaEventDestroy(start);
cudaEventDestroy(stop);

让我们编译并使用以下命令查看输出：

$ nvcc -m64 -run -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o cuda_event ./cuda_event.cu
compared a sample result...
host: 1.523750, device: 1.523750
CUDA event estimated - elapsed 23.408 ms 
Host measured time= 35.063 msec/s

如您所见，我们可以使用 CUDA 事件来测量内核执行时间。但是，测量的时间在 CUDA 事件和计时器之间存在间隙。我们可以使用 NVIDIA 分析器来验证哪个提供更准确的信息。当我们使用# nvprof ./cuda_event命令时，输出如下：

如您所见，与从主机测量相比，CUDA 事件提供了准确的结果。

使用 CUDA 事件的另一个好处是，我们可以使用多个 CUDA 流同时测量多个内核执行时间。让我们实现一个示例应用程序并查看其操作。

多流估计

cudaEventRecord()函数对主机是异步的。换句话说，没有同步来测量内核执行时间到示例代码。为了使事件和主机同步，我们需要使用cudaEventSynchronize()。例如，当我们在cudaEventRecord(stop)之后立即放置这个函数时，可以在设备到主机的异步数据传输之前放置内核函数打印，通过同步效果来实现。

在多个 CUDA 流应用程序中测量内核执行时间也是有用的：

让我们将这应用到04_stream_priority示例代码中的多个 CUDA 流重叠的代码中。使用以下代码更新代码：

class Operator
{
private:
    int _index;
    cudaStream_t stream;
    StopWatchInterface *p_timer;
    cudaEvent_t start, stop;

public:
    Operator() {
        cudaStreamCreate(&stream);

 // create cuda event
 cudaEventCreate(&start);
 cudaEventCreate(&stop);
    }

    ~Operator() {
        cudaStreamDestroy(stream);

 // destroy cuda event
 cudaEventDestroy(start);
 cudaEventDestroy(stop);
    }

    void set_index(int idx) { index = idx; }
    void async_operation(float *h_c, const float *h_a, 
                          const float *h_b,
                          float *d_c, float *d_a, float *d_b,
                          const int size, const int bufsize);
 void print_kernel_time();

}; // Operator

然后，我们将定义此时包含的print_time()函数，如下所示：

void Operator::print_time() {
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);
    printf("Stream %d time: %.4f ms\n", index, milliseconds);
}

现在，在Operator::async_operation()的开头和结尾插入cudaEventRecord()函数调用，如下所示：

void Operator::async_operation( ... )
{
    // start timer
    sdkStartTimer(&p_timer);

    // copy host -> device
    cudaMemcpyAsync(d_a, h_a, bufsize, 
                    cudaMemcpyHostToDevice, stream);
    cudaMemcpyAsync(d_b, h_b, bufsize, 
                    cudaMemcpyHostToDevice, stream);

    // record the event before the kernel execution
 cudaEventRecord(start, stream);

    // launch cuda kernel
    dim3 dimBlock(256);
    dim3 dimGrid(size / dimBlock.x);
    vecAdd_kernel<<< dimGrid, dimBlock, 0, 
                     stream >>>(d_c, d_a, d_b);

    // record the event right after the kernel execution finished
 cudaEventRecord(stop, stream);

    // copy device -> host
    cudaMemcpyAsync(h_c, d_c, bufsize, 
                    cudaMemcpyDeviceToHost, stream);

    // what happen if we include CUDA event synchronize?
    // QUIZ: cudaEventSynchronize(stop);

    // register callback function
    cudaStreamAddCallback(stream, Operator::Callback, this, 0);
}

对于这个函数，在函数的末尾放置同步是一个挑战。在完成本节后尝试这样做。这将影响应用程序的行为。建议尝试自己解释输出，然后使用分析器进行确认。

现在，让我们编译并查看执行时间报告，如下；它显示与先前执行类似的性能：

$ nvcc -m64 -run -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o cuda_event_with_streams ./cuda_event_with_streams.cu
Priority Range: low(0), high(-1)
stream 0 - elapsed 11.348 ms 
stream 3 - elapsed 19.435 ms 
stream 1 - elapsed 22.707 ms 
stream 2 - elapsed 35.768 ms 
kernel in stream 0 - elapsed 6.052 ms 
kernel in stream 1 - elapsed 14.820 ms 
kernel in stream 2 - elapsed 17.461 ms 
kernel in stream 3 - elapsed 6.190 ms 
compared a sample result...
host: 1.523750, device: 1.523750
Time= 35.993 msec, bandwidth= 22.373972 GB/s

在这个输出中，我们还可以看到每个内核的执行时间，这要归功于 CUDA 事件。从这个结果中，我们可以看到内核执行时间延长了，就像我们在上一节中看到的那样。

如果您想了解更多关于 CUDA 事件特性的信息，请查看 NVIDIA 的 CUDA 事件文档：docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__EVENT.html。

现在，我们将介绍管理 CUDA 网格的其他一些方面。第一项是动态并行性，它使 GPU 内核函数能够进行内核调用。

CUDA 动态并行性

CUDA 动态并行性（CDP）是一种设备运行时功能，它允许从设备函数进行嵌套调用。这些嵌套调用允许子网格具有不同的并行性。当问题需要不同的块大小时，此功能非常有用。

理解动态并行性

与主机的普通内核调用一样，GPU 内核调用也可以进行内核调用。以下示例代码显示了它的工作原理：

__global__ void child_kernel(int *data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    atomicAdd(&data[idx], seed);
}

__global__ void parent_kernel(int *data)
{
 if (threadIdx.x == 0) {
        int child_size = BUF_SIZE/gridDim.x;
        child_kernel<<< child_size/BLOCKDIM, BLOCKDIM >>>
                        (&data[child_size*blockIdx.x], blockIdx.x+1);
    }
    // synchronization for other parent's kernel output
    cudaDeviceSynchronize();
}

如您在这些函数中所见，我们需要确保哪个 CUDA 线程进行内核调用以控制网格创建的数量。要了解更多信息，让我们使用这个实现第一个应用程序。

动态并行性的使用

我们的动态并行性代码将创建一个父网格，该父网格将创建一些子网格：

首先，我们将使用以下代码编写parent_kernel()函数和child_kernel()函数：

#define BUF_SIZE (1 << 10)
#define BLOCKDIM 256

__global__ void child_kernel(int *data)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    atomicAdd(&data[idx], 1);
}

__global__ void parent_kernel(int *data)
{
    if (blockIdx.x * blockDim.x + threadIdx.x == 0)
    {
        int child_size = BUF_SIZE/gridDim.x;
        child_kernel<<< child_size/BLOCKDIM, BLOCKDIM >>> \
                        (&data[child_size*blockIdx.x], 
                         blockIdx.x+1);
    }
    // synchronization for other parent's kernel output
    cudaDeviceSynchronize();
}

如您在这段代码中所见，父内核函数创建子内核网格作为块的数量。然后，子网格递增指定的内存1来标记它们的操作。内核执行后，父内核使用cudaDeviceSynchronize()函数等待所有子网格完成其工作。在进行同步时，我们应确定同步的范围。如果我们需要在块级别进行同步，我们应选择__synchthread()。

使用以下代码编写main()函数：

#define BUF_SIZE (1 << 10)
#define BLOCKDIM 256
int main()
{
    int *data;
    int num_child = 4;

    cudaMallocManaged((void**)&data, BUF_SIZE * sizeof(int));
    cudaMemset(data, 0, BUF_SIZE * sizeof(int));

    parent_kernel<<<num_child, 1>>>(data);
    cudaDeviceSynchronize();

    // Count elements value
    int counter = 0;
    for (int i = 0; i < BUF_SIZE; i++)
        counter += data[i];

    // getting answer
    int counter_h = 0;
    for (int i = 0; i < num_child; i++)
        counter_h += (i+1);
    counter_h *= BUF_SIZE / num_child;

    if (counter_h == counter)
        printf("Correct!!\n");
    else
        printf("Error!! Obtained %d. It should be %d\n", 
               counter, counter_h);

    cudaFree(data);
    return 0;
}

正如前面讨论的，我们将创建子网格以及块的数量。因此，我们将使用网格大小为4来执行父内核函数，而块大小为1。

要编译 CDP 应用程序，我们应该为nvcc编译器提供-rdc=true选项。因此，编译源代码的命令如下：

$ nvcc -run -rdc=true -lcudadevrt -gencode arch=compute_70,code=sm_70 -o host_callback host_callback.cu -I/usr/local/cuda/samples/common/inc

让我们对这个应用程序进行分析，以了解其操作。以下截图显示了这个嵌套调用的工作原理：

如我们在这个屏幕截图中所见，父内核创建了一个子网格，我们可以在左侧面板的右角标中看到它们的关系。然后，父网格（parent_kernel）等待其执行，直到子网格完成其工作。CUDA 目前不支持 SM70（Volta 架构）的 CDT 分析，因此我使用 Tesla P40 来获得这个输出。

递归

动态并行性的一个好处是我们可以创建递归。以下代码显示了一个递归内核函数的示例：

__global__ void recursive_kernel(int *data, int size, int depth) {
  int x_0 = blockIdx.x * size;

  if (depth > 0) {
    __syncthreads();
 if (threadIdx.x == 0) {
        int dimGrid = size / dimBlock;
        recursive_kernel<<<dimGrid, 
              dimBlock>>>(&data[x_0], size/dimGrid, depth-1);
        cudaDeviceSynchronize();
      }
      __syncthreads();
   }
}

如您所见，与以前的动态并行内核函数相比，没有太大的区别。但是，我们应该谨慎使用这个功能，考虑到资源使用和限制。一般来说，动态并行内核可以保守地保留高达 150MB 的设备内存来跟踪待处理的网格启动和通过在子网格启动上进行同步来同步父网格的状态。此外，同步必须在多个级别上小心进行，而嵌套内核启动的深度限制为 24 级。最后，控制嵌套内核启动的运行时可能会影响整体性能。

如果您需要了解动态并行性的限制和限制，请参阅以下编程指南：docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#implementation-restrictions-and-limitations。

我们将在第七章中介绍其在 CUDA 中快速排序实现中的应用，即CUDA 中的并行编程模式。要了解更多关于动态并行性的信息，请参阅以下文档：

网格级别的协作组

如第三章中所讨论的，CUDA 提供了协作组。协作组可以根据其分组目标进行分类：warp 级别、块级别和网格级别的组。本文介绍了网格级别的协作组，并探讨了协作组如何处理 CUDA 网格。

协作组最显著的好处是对目标并行对象的显式同步。使用协作组，程序员可以设计他们的应用程序来显式同步 CUDA 并行对象、线程块或网格。使用第三章中介绍的块级协作组，CUDA 线程编程，我们可以通过指定需要同步的 CUDA 线程或块来编写更易读的代码。

理解网格级协作组

自 9.0 版本以来，CUDA 提供了另一级协作组，与网格一起工作。具体来说，有两个网格级协作组：grid_group和multi_grid_group。使用这些组，程序员可以描述网格在单个 GPU 或多个 GPU 上的操作同步。

在这个示例中，我们将探索grid_group的功能，它可以同步网格与减少问题，就像第三章中所提到的，CUDA 线程编程，关于基于块级减少的先前减少设计。每个线程块产生自己的减少结果，并将它们存储到全局内存中。然后，另一个块级减少内核启动，直到我们获得单个减少值。这是因为完成内核操作可以保证下一个减少内核从多个线程块中读取减少值。其设计由左侧的图表描述：

！

另一方面，网格级同步使另一种内部同步块式减少结果的内核设计成为可能，以便主机只需调用一次内核即可获得减少结果。在协作组中，grid_group.sync()提供了这样的功能，因此我们可以编写减少内核而无需内核级迭代。

要使用grid_group.sync()函数，我们需要使用cudaLaunchCooperativeKernel()函数调用内核函数。其接口设计如下：

__host__ cudaError_t cudaLaunchCooperativeKernel
    ( const T* func, dim3 gridDim, dim3 blockDim, 
      void** args, size_t sharedMem = 0, cudaStream_t stream = 0 )

因此，它的使用方式与cudaLaunchKernel()函数相同，该函数启动内核函数。

为了使grid_group中的所有线程块同步，网格中活动线程块的总数不应超过内核函数和设备的最大活动块数。GPU 上的最大活动块大小是每个 SM 的最大活动块数和流处理器的数量的乘积。违反此规则可能导致死锁或未定义行为。我们可以使用cudaOccupancyMaxActiveBlocksPerMultiprocessor()函数来获取每个 SM 内核函数的最大活动线程块数，通过传递内核函数和块大小信息。

使用`grid_group`的用法

现在，让我们将grid_group应用于并行减少问题，并看看 GPU 编程如何改变：

我们将重用之前并行减少代码中的主机代码，即03_cuda_thread_programming/07_cooperative_groups。换句话说，我们将通过对主机代码进行小的更改来改变 GPU 的操作。您还可以使用07_grid_level_cg目录中的代码。
现在，让我们编写一些块级减少代码。当我们有网格级协作组时，所有线程块必须是活动的。换句话说，我们不能执行多个线程块，而 GPU 能够执行的活动块。因此，这个减少将首先累积输入数据，以覆盖所有数据，使用有限数量的线程块。然后，它将在块级进行并行减少，就像我们在第三章中所介绍的那样，CUDA 线程编程。

以下代码显示了它的实现：

__device__ void
block_reduction(float *out, float *in, float *s_data, int active_size, int size, 
          const cg::grid_group &grid, const cg::thread_block &block)
{
  int tid = block.thread_rank();

  // Stride over grid and add the values to a shared memory buffer
  s_data[tid] = 0.f;
  for (int i = grid.thread_rank(); i < size; i += active_size)
    s_data[tid] += in[i];

  block.sync();

  for (unsigned int stride = blockDim.x / 2; 
       stride > 0; stride >>= 1) {
    if (tid < stride)
      s_data[tid] += s_data[tid + stride];
    block.sync();
  }

  if (block.thread_rank() == 0)
    out[block.group_index().x] = s_data[0];
}

然后，让我们编写一个内核函数，考虑活动块数和grid_group执行块级减少。在这个函数中，我们将调用块级减少代码，并在网格级别进行同步。然后，我们将从输出中执行并行减少，就像我们在第三章 CUDA 线程编程中所介绍的那样。以下代码显示了其实现：

__global__ void
reduction_kernel(float *g_out, float *g_in, unsigned int size)
{
  cg::thread_block block = cg::this_thread_block();
  cg::grid_group grid = cg::this_grid();
  extern __shared__ float s_data[];

  // do reduction for multiple blocks
  block_reduction(g_out, g_in, s_data, grid.size(), 
                  size, grid, block);

  grid.sync();

  // do reduction with single block
  if (block.group_index().x == 0)
    block_reduction(g_out, g_out, s_data, block.size(), gridDim.x, grid, block);
}

最后，我们将实现调用具有可用活动线程块维度的内核函数的主机代码。为此，此函数使用cudaoccupancyMaxActiveBlocksPerMultiprocessor()函数。此外，网格级合作组要求我们通过cudaLaunchCooperativeKernel()函数调用内核函数。您可以在这里看到实现：

int reduction_grid_sync(float *g_outPtr, float *g_inPtr, int size, int n_threads)
{ 
  int num_blocks_per_sm;
  cudaDeviceProp deviceProp;

  // Calculate the device occupancy to know 
  // how many blocks can be run concurrently
  cudaGetDeviceProperties(&deviceProp, 0);
  cudaOccupancyMaxActiveBlocksPerMultiprocessor(&num_blocks_per_sm, 
      reduction_kernel, n_threads, n_threads*sizeof(float));
  int num_sms = deviceProp.multiProcessorCount;
  int n_blocks = min(num_blocks_per_sm * num_sms, 
                     (size + n_threads - 1) / n_threads);

  void *params[3];
  params[0] = (void*)&g_outPtr;
  params[1] = (void*)&g_inPtr;
  params[2] = (void*)&size;
  cudaLaunchCooperativeKernel((void*)reduction_kernel, 
                              n_blocks, n_threads, params, 
                              n_threads * sizeof(float), NULL);

  return n_blocks;
}

现在，请确保可以从reduction.cpp文件中调用主机函数。
然后，让我们编译代码并查看其操作。以下 shell 命令编译代码并执行应用程序。计算能力应该等于或大于70：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -rdc=true -o reduction ./reduction.cpp ./reduction_kernel.cu
Time= 0.474 msec, bandwidth= 141.541077 GB/s
host: 0.996007, device 0.996007

输出性能远远落后于我们在第三章 CUDA 线程编程的最终结果。由于block_reduction()函数在开始时使用了高内存吞吐量，因此它是高度内存绑定的：

主要影响因素是我们只能使用活动线程块。因此，我们无法隐藏内存访问时间。实际上，使用grid_group还有其他目的，例如图搜索、遗传算法和粒子模拟，这要求我们保持状态长时间处于活动状态以获得性能。

这种网格级同步可以为性能和可编程性提供更多好处。由于这使得内核可以自行同步，我们可以使内核自行迭代。因此，它对解决图搜索、遗传算法和实际模拟非常有用。要了解有关grid_groups中合作组的更多信息，请参阅提供的文档on-demand.gputechconf.com/gtc/2017/presentation/s7622-Kyrylo-perelygin-robust-and-scalable-cuda.pdf。

使用 OpenMP 的 CUDA 内核调用

为了增加应用程序的并发性，我们可以从主机的并行任务中进行内核调用。例如，OpenMP 提供了多核架构的简单并行性。本教程介绍了 CUDA 如何操作 OpenMP。

OpenMP 和 CUDA 调用

OpenMP 使用分叉-合并模型的并行性来针对多核 CPU。主线程启动并行操作并创建工作线程。主机线程并行运行自己的工作，并在完成工作后加入。

使用 OpenMP，CUDA 内核调用可以与多个线程并行执行。这有助于程序员不必维护单独的内核调用，而是允许它们的内核执行依赖于主机线程的索引。

在本节中，我们将使用以下 OpenMP API：

omp_set_num_threads()设置将并行工作的工作线程数。
omp_get_thread_num()返回工作线程的索引，以便每个线程可以识别其任务。
#pragma omp parallel {} 指定了一个并行区域，将由工作线程覆盖。

现在，让我们编写一些代码，其中 OpenMP 调用 CUDA 内核函数。

CUDA 与 OpenMP 的内核调用

在本节中，我们将实现一个使用 OpenMP 的多流矢量加法应用程序。为此，我们将修改先前的版本并查看差异：

要测试 CUDA 中的 OpenMP，我们将修改03_cuda_callback目录中的代码。我们将修改main()函数的主体，或者您可以使用放置在08_openmp_cuda目录中的提供的示例代码。
现在，让我们包括 OpenMP 头文件并修改代码。要在代码中使用 OpenMP，我们应该使用#include <omp.h>。而且，我们将更新代码，使其使用 OpenMP 来迭代每个流：

// execute each operator collesponding data
omp_set_num_threads(num_operator);
#pragma omp parallel
{
    int i = omp_get_thread_num();
    printf("Launched GPU task %d\n", i);

    int offset = i * size / num_operator;
    ls_operator[i].set_index(i);
    ls_operator[i].async_operation(&h_c[offset], &h_a[offset],   
                                   &h_b[offset],&d_c[offset], 
                                   &d_a[offset], &d_b[offset],
                                   size / num_operator, bufsize 
                                   / num_operator);
}

使用以下命令编译代码：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -Xcompiler -fopenmp -lgomp -o openmp ./openmp.cu
stream 0 - elapsed 10.734 ms 
stream 2 - elapsed 16.153 ms 
stream 3 - elapsed 21.968 ms 
stream 1 - elapsed 27.668 ms 
compared a sample result...
host: 1.523750, device: 1.523750
Time= 27.836 msec, bandwidth= 28.930389 GB/s

每当您执行此应用程序时，您将看到每个流以无序方式完成其工作。此外，每个流显示不同的时间。这是因为 OpenMP 可以创建多个线程，并且操作是在运行时确定的。

为了了解其运行情况，让我们对应用程序进行分析。以下截图显示了应用程序的分析时间表。由于调度的原因，这可能与您的情况不同：

如您在此截图中所见，您将能够看到数据传输与 Stream 17 相比已经反转。因此，我们可以看到第二个流最终完成了它的工作。

多进程服务

GPU 能够从并发的 CPU 进程中执行内核。但是，默认情况下，它们只以分时方式执行，即使每个内核没有充分利用 GPU 计算资源。为了解决这种不必要的串行化，GPU 提供了多进程服务（MPS）模式。这使得不同的进程能够同时在 GPU 上执行它们的内核，以充分利用 GPU 资源。启用时，nvidia-cuda-mps-control守护进程监视目标 GPU，并使用该 GPU 管理进程内核操作。此功能仅在 Linux 上可用。在这里，我们可以看到多个进程共享同一个 GPU 的 MPS：

正如我们所看到的，每个进程在 GPU 上并行运行一部分（绿色条），而一部分在 CPU 上运行（蓝色条）。理想情况下，您需要蓝色条和绿色条都能获得最佳性能。这可以通过利用所有最新 GPU 支持的 MPS 功能来实现。

请注意，当一个 MPI 进程无法饱和整个 GPU 并且代码的重要部分也在 CPU 上运行时，多个 MPI 进程在同一个 GPU 上运行是有益的。如果一个 MPI 进程利用整个 GPU，即使 CPU 部分（蓝色条）会减少，绿色条的时间也不会减少，因为 GPU 完全被一个 MPI 进程利用。其他 MPI 进程将根据 GPU 架构以分时方式依次访问 GPU。这类似于启动并发内核的情况。如果一个内核利用整个 GPU，那么另一个内核要么等待第一个内核完成，要么进行分时。

这样做的好处是不需要对应用程序进行任何更改即可使用 MPS。MPS 进程作为守护进程运行，如下命令所示：

$nvidia-smi -c EXCLUSIVE_PROCESS 
$nvidia-cuda-mps-control –d

运行此命令后，所有进程都将其命令提交给 MPS 守护进程，该守护进程负责将 CUDA 命令提交给 GPU。对于 GPU，只有一个进程访问 GPU（MPS 守护进程），因此多个进程可以同时运行来自多个进程的多个内核。这可以帮助将一个进程的内存复制与其他 MPI 进程的内核执行重叠。

消息传递接口简介

消息传递接口（MPI）是一种并行计算接口，它能够触发多个进程跨计算单元 - CPU 核心、GPU 和节点。典型的密集多 GPU 系统包含 4-16 个 GPU，而 CPU 核心的数量在 20-40 个之间。在启用 MPI 的代码中，应用程序的某些部分作为不同的 MPI 进程在多个核心上并行运行。每个 MPI 进程都将调用 CUDA。了解将 MPI 进程映射到相应的 GPU 非常重要。最简单的映射是 1:1，即每个 MPI 进程都独占相应的 GPU。此外，我们还可以将多个 MPI 进程理想地映射到单个 GPU 上。

为了将多进程应用场景应用到单个 GPU 上，我们将使用 MPI。要使用 MPI，您需要为您的系统安装 OpenMPI。按照以下步骤在 Linux 上安装 OpenMPI。此操作已在 Ubuntu 18.04 上进行了测试，因此如果您使用其他发行版，可能会有所不同：

$ wget -O /tmp/openmpi-3.0.4.tar.gz https://www.open-mpi.org/software/ompi/v3.0/downloads/openmpi-3.0.4.tar.gz
$ tar xzf /tmp/openmpi-3.0.4.tar.gz -C /tmp
$ cd /tmp/openmpi-3.0.4
$ ./configure --enable-orterun-prefix-by-default --with-cuda=/usr/local/cuda
$ make -j $(nproc) all && sudo make install
$ sudo ldconfig
$ mpirun --version
mpirun (Open MPI) 3.0.4

Report bugs to http://www.open-mpi.org/community/help/

现在，让我们实现一个可以与 MPI 和 CUDA 一起工作的应用程序。

实现一个启用 MPI 的应用程序

要使应用程序与 MPI 一起工作，我们需要在应用程序中放入一些可以理解 MPI 命令的代码：

我们将重用 OpenMP 示例代码，因此将openmp.cu文件复制到08_openmp_cuda目录中。
在代码开头插入mpi头文件include语句：

#include <mpi.h>

在main()函数中创建秒表后立即插入以下代码：

// set num_operator as the number of requested process
int np, rank;
MPI_Init(&argc, &argv);
MPI_Comm_size(MPI_COMM_WORLD, &np);
MPI_Comm_rank(MPI_COMM_WORLD, &rank);

按照第 3 步中提到的代码后，将所需的内存大小切割为进程数，如下所示：

bufsize /= np;
size /= np;

我们需要让每个线程报告它们所属的进程。让我们更新并行执行代码块中的printf()函数，如下所示：

// execute each operator collesponding data
omp_set_num_threads(num_operator);
#pragma omp parallel
{
    int i = omp_get_thread_num();
    int offset = i * size / num_operator;
    printf("Launched GPU task (%d, %d)\n", rank, i);

    ls_operator[i].set_index(i);
    ls_operator[i].async_operation(&h_c[offset], 
                                   &h_a[offset], &h_b[offset],
                                   &d_c[offset], &d_a[offset], 
                                   &d_b[offset],
                                   size / num_operator, 
                                   bufsize / num_operator);
}

在main()的末尾放置MPI_Finalize()函数以关闭 MPI 实例。
使用以下命令编译代码：

$ nvcc -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -I/usr/local/include/ -Xcompiler -fopenmp -lgomp -lmpi -o simpleMPI ./simpleMPI.cu

您必须使用 GPU 的计算能力版本号来选择gencode选项。

使用以下命令测试编译后的应用程序：

$ ./simpleMPI 2

现在，使用以下命令测试 MPI 执行：

$ mpirun -np 2 ./simpleMPI 2
Number of process: 2
Number of operations: 2
Launched GPU task (1, 0)
Launched GPU task (1, 1)
Number of operations: 2
Launched GPU task (0, 0)
Launched GPU task (0, 1)
stream 0 - elapsed 13.390 ms 
stream 1 - elapsed 25.532 ms 
compared a sample result...
host: 1.306925, device: 1.306925
Time= 25.749 msec, bandwidth= 15.637624 GB/s
stream 0 - elapsed 21.334 ms 
stream 1 - elapsed 26.010 ms 
compared a sample result...
host: 1.306925, device: 1.306925
Time= 26.111 msec, bandwidth= 15.420826 GB/s

启用 MPS

在 GPU 上启用 MPS 需要对 GPU 操作模式进行一些修改。但是，您需要具有比 Kepler 架构更晚的 GPU 架构。

让我们按照以下步骤启用 MPS：

使用以下命令启用 MPS 模式：

$ export CUDA_VISIBLE_DEVICES=0
$ sudo nvidia-smi -i 0 -c 3
$ sudo nvidia-cuda-mps-control -d

或者，您可以使用make enable_mps命令来使用此预定义在Makefile中的配方示例代码。然后，我们可以从nivida-smi输出中看到更新后的计算模式：

现在，使用以下命令测试 MPS 模式下的 MPI 执行：

$ mpirun -np 2 ./simpleMPI 2
Number of process: 2
Number of operations: 2
Launched GPU task (1, 0)
Launched GPU task (1, 1)
stream 0 - elapsed 10.203 ms 
stream 1 - elapsed 15.903 ms 
compared a sample result...
host: 1.306925, device: 1.306925
Time= 16.129 msec, bandwidth= 24.964548 GB/s
Number of operations: 2
Launched GPU task (0, 0)
Launched GPU task (0, 1)
stream 0 - elapsed 10.203 ms 
stream 1 - elapsed 15.877 ms 
compared a sample result...
host: 1.306925, device: 1.306925
Time= 15.997 msec, bandwidth= 25.170544 GB/s

如您所见，与之前的执行相比，每个进程的经过时间都有所减少。

现在，让我们恢复原始模式。要禁用 MPS 模式，请使用以下命令：

$ echo "quit" | sudo nvidia-cuda-mps-control
$ sudo nvidia-smi -i 0 -c 0

或者，您可以使用make disable_mps命令来使用此预定义在Makefile中的配方示例代码。

要了解更多关于 MPS 的信息，请使用以下链接：

对 MPI 应用程序进行分析并了解 MPS 操作

使用 MPI，多个进程的内核可以同时共享 GPU 资源，从而增强整体 GPU 利用率。没有 MPS，由于时间切片共享和上下文切换开销，GPU 资源被低效地共享。

以下屏幕截图显示了没有 MPS 的多个进程的时间轴配置文件结果：

在此配置文件中，我们可以看到两个 CUDA 上下文共享一个 GPU，并且由于上下文之间的时间共享，内核执行时间延长。

另一方面，MPS 模式管理内核执行请求，因此所有内核执行都会像使用单个进程一样启动。以下屏幕截图显示了 MPS 模式下的内核执行：

如您所见，只有一个 CUDA 流驻留在 GPU 上并控制所有 CUDA 流。此外，所有内核执行时间都得到了稳定，并且使用 MPS 可以减少总的经过时间。总之，使用 MPS 模式有利于多个 GPU 进程的整体性能，并共享 GPU 资源。

nvprof支持将多个 MPI 进程的分析器信息转储到不同的文件中。例如，对于基于 Open MPI 的应用程序，以下命令将在多个文件中转储分析信息，每个文件的名称都基于 MPI 进程的排名：

$ mpirun -np 2 nvprof -f -o simpleMPI.%q{OMPPI_COMM_WORLD_RANK}_2.nvvp ./simpleMPI 2

或者，您可以使用以下命令来执行示例代码：

$ PROCS=2 STREAMS=2 make nvprof

然后，您将为每个进程获得两个nvvp文件。

现在，我们将使用以下步骤使用 NVIDIA Visual Profiler 来查看这些nvvp文件：

打开文件|导入菜单，通过导入nvvp文件创建一个分析会话：

在 Windows 或 Linux 中，快捷键是Ctrl + I，OSX 使用command + I。

然后从列表中选择 Nvprof 后，点击下一步按钮：

从 Nvprof 选项中，选择多个进程，然后单击下一步>：

从导入 Nvprof 数据中，单击浏览...按钮，并选择由nvprof生成的nvvp文件。要对具有多个进程的应用程序进行分析，您需要导入nvvp文件，因为存在多个进程：

单击完成，然后 NVIDIA Visual Profiler 将以时间线视图显示分析结果，如下所示：

请注意，只有同步 MPI 调用将由nvprof进行注释。如果使用异步 MPI API，则需要使用其他 MPI 专用的分析工具。其中一些最著名的工具包括以下内容：

TAU：TAU 是一种性能分析工具包，目前由俄勒冈大学维护。
Vampir：这是一种商业可用的工具，对数百个 MPI 进程具有良好的可伸缩性。
Intel VTune Amplifier：商业工具的另一个选择是 Intel VTune Amplifier。它是目前可用的最好的工具之一，可用于 MPI 应用程序分析。

最新的 CUDA 工具包还允许对 MPI API 进行注释。为此，需要将--annotate-mpi标志传递给nvprof，如以下命令所示：

mpirun -np 2 nvprof --annotate-mpi openmpi -o myMPIApp.%q{OMPI_COMM_WORLD_RANK}.nvprof ./myMPIApplciation

内核执行开销比较

对于迭代并行 GPU 任务，我们有三种内核执行方法：迭代内核调用，具有内部循环，以及使用动态并行性进行递归。最佳操作由算法和应用程序确定。但是，您也可以考虑它们之间的内核执行选项。本示例帮助您比较这些内核执行开销并审查它们的可编程性。

首先，让我们确定我们将测试哪种操作。本示例将使用一个简单的 SAXPY 操作。这有助于我们专注并制作迭代执行代码。此外，随着操作变得更简单，操作控制开销将变得更重。但是，您当然可以尝试任何其他操作。

实现三种内核执行方式

以下步骤涵盖了三种不同迭代操作的性能比较：

创建并导航到10_kernel_execution_overhead目录。
编写simple_saxpy_kernel()函数，代码如下：

__global__ void
simple_saxpy_kernel(float *y, const float* x, const float alpha, const float beta)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    y[idx] = alpha * x[idx] + beta;
}

编写iterative_saxpy_kernel()函数，代码如下：

__global__ void
iterative_saxpy_kernel(float *y, const float* x, 
                       const float alpha, const float beta, 
                       int n_loop)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    for (int i = 0; i < n_loop; i++)
        y[idx] = alpha * x[idx] + beta;
}

编写recursive_saxpy_kernel()函数，代码如下：

__global__ void
recursive_saxpy_kernel(float *y, const float* x, 
                       const float alpha, const float beta, 
                       int depth)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    if (depth == 0)
        return;
    else
        y[idx] = alpha * x[idx] + beta;

    if (idx == 0)
        vecAdd_kernel_C<<< gridDim.x, blockDim.x 
                           >>>(y, x, alpha, beta, depth - 1);
}

编写启动这些 CUDA 内核函数的主机代码。首先，我们将对simple_saxpy_kernel()函数进行迭代调用：

for (int i = 0; i < n_loop; i++) {
    simple_saxpy_kernel<<< dimGrid, dimBlock >>>(
                           d_y, d_x, alpha, beta);
}

其次，我们将调用iterative_saxpy_kernel()内核函数，该函数内部有一个迭代循环：

iterative_saxpy_kernel<<< dimGrid, dimBlock >>>(
                          d_y, d_x, alpha, beta, n_loop);

最后，我们将调用recursive_saxpy_kernel()内核函数，该函数以递归方式调用自身：

recursive_saxpy_kernel<<< dimGrid, dimBlock >>>(
                          d_y, d_x, alpha, beta, n_loop);

循环次数小于或等于 24，因为最大递归深度为 24。除了简单的循环操作外，您不必在主机上放置循环操作，因为它已在内核代码中定义。

使用以下命令编译代码：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -rdc=true -o cuda_kernel ./cuda_kernel.cu

您必须使用 GPU 的计算能力版本号来选择gencode选项。

测试编译后的应用程序。这个结果是使用 Tesla P40 测量的，因为 CUDA 9.x 不支持 Volta GPU 的 CUDA 动态并行性（CDP）配置文件：

Elapsed Time...
simple loop: 0.094 ms
inner loop : 0.012 ms
recursion : 0.730 ms

三种执行的比较

从结果中，我们可以确认内部循环是迭代操作中最快的方法。以下截图显示了这个示例应用程序的分析结果：

迭代内核调用显示了每个内核调用的内核启动开销。GPU 需要从设备内存中获取所有所需的数据，并需要调度 GPU 资源等。另一方面，内部循环内核显示了一个打包操作，因为所有所需的资源都是预先定位的，不需要重新调度其执行。由于我们之前讨论的动态并行性限制，递归内核操作显示了最长的执行时间。

一般来说，建议使用开销最小的方法。然而，很难说哪种内核调用设计优于其他，因为算法和问题比我们在这里涵盖的要多。例如，CDP 用于增强某些情况下的并行性，比如用于 GPU 树和搜索。

总结

在本章中，我们涵盖了几种内核执行机制。我们讨论了 CUDA 流是什么，以及如何使用它们同时执行多个内核函数。通过利用主机和 GPU 之间的异步操作，我们学到可以通过数据传输和内核执行来隐藏内核执行时间。此外，我们可以使用回调函数使 CUDA 流调用主机函数。我们可以创建一个有优先级的流，并确认其有优先级的执行。为了测量内核函数的确切执行时间，我们使用了 CUDA 事件，并且我们也学到 CUDA 事件可以用于与主机同步。在最后一节中，我们还讨论了每种内核执行方法的性能。

我们还涵盖了其他内核操作模型：动态并行性和网格级协作组。动态并行性使得内核函数内部可以进行内核调用，因此我们可以使用递归操作。网格级协作组实现了多功能的网格级同步，我们讨论了这个特性在特定领域的用途：图搜索、遗传算法和粒子模拟。

然后，我们扩展了我们对主机的覆盖范围。CUDA 内核可以从多个线程或多个进程中调用。为了执行多个线程，我们使用了带有 CUDA 的 OpenMP，并讨论了它的用处。我们使用 MPI 来模拟多进程操作，并且可以看到 MPS 如何提高整体应用程序性能。

正如我们在本章中看到的，选择正确的内核执行模型是一个重要的话题，线程编程也是如此。这可以优化应用程序的执行时间。现在，我们将扩展我们的讨论到多 GPU 编程来解决大问题。

第五章：CUDA 应用程序分析和调试

CUDA 为开发人员提供了许多编程工具。这些工具包括编译器、分析器、IDE 及其插件、调试器和内存检查器。了解这些工具将有助于您分析您的应用程序，并帮助您完成我们将要涵盖的开发项目。在本章中，我们将介绍这些工具的基本用法，并讨论如何将它们应用到应用程序开发中。

本章将涵盖以下主题：

在 GPU 应用程序中进行专注的分析目标范围
针对远程机器的可视化分析
使用 CUDA 错误调试 CUDA 应用程序
使用 CUDA Assert 断言本地 GPU 值
使用 Nsight Visual Studio Edition 调试 CUDA 应用程序
使用 Nsight Eclipse Edition 调试 CUDA 应用程序
使用 CUDA-GDB 调试 CUDA 应用程序
使用 CUDA-memcheck 进行运行时验证

技术要求

为了完成本章，建议您使用 Pascal 架构之后的 NVIDIA GPU 卡。换句话说，您的 GPU 的计算能力应该等于或大于 60。如果您不确定您的 GPU 架构，请访问 NVIDIA 的网站developer.nvidia.com/cuda-gpus，并确认您的 GPU 的计算能力。

本章的示例代码已经使用 CUDA Toolkit 的 10.1 版本进行开发和测试。一般来说，如果适用的话，建议您使用最新的 CUDA 版本。

在 GPU 应用程序中进行专注的分析目标范围

NVIDIA 的 Visual Profiler 是一个方便的工具，用于找出 GPU 应用程序中的瓶颈并理解它们的操作。虽然它提供了应用程序操作的流畅信息，但如果您只想专注于特定代码区域，这些信息可能会显得多余。在这种情况下，限制分析范围更加高效。

分析目标可以是特定的代码块、GPU 和时间。指定代码块称为专注分析。当您想要专注于特定内核函数的分析，或者在大型 GPU 应用程序的一部分上进行分析时，这种技术是有用的。在我们介绍专注分析后，将介绍针对 GPU 或时间的分析目标。

限制代码中的分析目标

为了从专注的分析中受益，您可能希望在源代码中包含特色的头文件，如下所示：

#include <cuda_profiler_api.h>

然后，您可以使用cudaProfilerStart()和cudaProfilerStop()来指定您的分析范围：

cudaProfilerStart();
... {target of profile} ...
cudaProfilerStop();

现在，您需要使用特定标志--profile-from-start来分析您的应用程序。

这个选项不会让分析器开始分析，直到请求到达。如果您想使用 NVIDIA Visual Profiler 来分析您的应用程序，请确保在设置视图中勾选“启动时启用分析”复选框。

以下步骤涵盖了如何使用一些简单的示例代码来控制 NVIDIA 分析器。为了使这更容易，我们将重用我们在第三章中用于矩阵乘法操作的示例代码，CUDA 线程编程：

编写一个 CUDA 应用程序，其中包含两个简单的 SGEMM CUDA 内核函数。这两个内核函数是相同的，但名称不同，即sgemm_kernel_A()和sgemm_kernel_B()。
进行两次迭代调用，如下所示：

int n_iter = 5;
for (int i = 0; i < n_iter; i++)
    sgemm_gpu_A(d_A, d_B, d_C, N, M, K, alpha, beta);
for (int i = 0; i < n_iter; i++)
    sgemm_gpu_B(d_A, d_B, d_C, N, M, K, alpha, beta);

现在，让我们编译代码并使用nvprof进行分析：

$ nvcc -m64 -gencode arch=compute_70,code=sm_70 -o sgemm sgemm.cu
$ nvprof -f -o profile-original.nvvp ./sgemm

当您使用 Visual Profiler 打开生成的profile-original.nvvp文件时，您将得到如下的分析结果：

这个时间轴包括了应用程序启动时的整个分析信息。然而，当我们想要优化我们的内核函数时，我们可以说分析结果包含了不必要的信息。

以下步骤涵盖了如何指定分析专注区域：

在源代码顶部放置 #include <cuda_profiler_api.h> 以启用专注分析 API。然后，我们可以使用 cudaProfilerStart() 和 cudaProfilerStop() 来包含我们感兴趣的区域，如下所示：

cudaProfilerStart();
for (int i = 0; i < n_iter; i++)
    sgemm_gpu_B(d_A, d_B, d_C, N, M, K, alpha, beta);
cudaProfilerStop();

编译您的代码并使用 Visual Profiler 查看更新后的分析结果。我们必须向分析器提供 --profile-from-start off 选项，如下所示：

$ nvcc -m64 -gencode arch=compute_70,code=sm_70 -o sgemm sgemm.cu
$ nvprof -f -o profile-start-stop.nvvp --profile-from-start off ./sgemm

当您打开新生成的分析结果时，分析器只会报告应用程序的指定部分，如下所示：

分析结果受限。上面的屏幕截图显示了内核执行的情况，从开始 GPU 执行时开始。因此，您可以省去对应用程序初始化和其他无关操作进行分析的步骤。

总之，专注分析有几个好处，如下所示：

这有助于您专注于当前正在开发的模块。
它可以让您在分析报告中删除无关的操作，例如：
与您的代码无关的外部模块行为
应用程序初始化延迟
在时间轴视图中查找目标函数时，这有助于节省时间。

通过时间或 GPU 限制分析目标

NVIDIA 分析器还有其他可以限制分析目标的选项。您也可以使用以下选项进行专注分析：

--timeout <second> 选项限制应用程序的执行时间。当您需要分析执行时间较长的迭代操作的应用程序时，此选项非常有用。
--devices <gpu ids> 选项指定要进行分析的 GPU。该选项帮助您在多 GPU 应用程序中缩小 GPU 内核操作的范围。

此外，如果您只想专注于少数内核函数，您不必收集所有指标。您可以使用 --kernels、--event 和 --metrics 选项向分析器指定您的兴趣。您可以将这些选项与其他分析选项一起使用，如下所示：

$ nvprof -f -o profile_kernels_metric.nvvp --kernels sgemm_kernel_B --metrics all ./sgemm

将收集的指标导入时间轴分析结果后，您会发现目标内核只有指标信息。

在 CPU 抽样中有许多其他多功能的分析特性，例如标记分析范围、OpenMP 和 OpenACC 分析等。如果您想了解 NVIDIA 分析器的功能，请查看 NVIDIA 的 Jeff Larkin 提供的以下分析器介绍讲座：www.olcf.ornl.gov/wp-content/uploads/2018/12/summit_workshop_Profilers.pdf。

NVIDIA 的官方分析器用户指南提供了有关 NVIDIA 分析器功能的详细信息 (docs.nvidia.com/cuda/profiler-users-guide/index.html).

使用 NVTX 进行分析

通过专注分析，我们可以使用 cudaProfilerStart() 和 cudaProfilerStop() 对有限的特定区域进行分析。但是，如果我们想要分析复杂应用程序中的功能性能，这是有限的。对于这种情况，CUDA 分析器通过 NVIDIA 工具扩展 (NVTX) 提供时间轴注释。

使用 NVTX，我们可以对 CUDA 代码进行注释。我们可以使用 NVTX API 如下：

nvtxRangePushA("Annotation");
.. { Range of GPU operations } ..
cudaDeviceSynchronization();     // in case if the target code block is pure kernel calls
nvtxRangePop();

如您所见，我们可以将一段代码定义为一组代码，并手动注释该范围。然后，CUDA 分析器提供注释的时间轴跟踪，以便我们可以测量代码块的执行时间。这种方法的一个缺点是 NVTX API 是主机函数，因此如果目标代码块是纯 GPU 内核调用，则需要同步主机和 GPU。

要了解更多信息，请将此 NVTX 代码应用于前面的专注分析示例。首先，我们应该包含一个 NVTX 头文件，如下所示：

#include "nvToolsExt.h"

然后，我们将在几个地方插入 nvtxRangePushA() 和 nvtxRangePop()，如下所示：

    cudaProfileStart();
    // copy initial value for gpu memory
    nvtxRangePushA("Data Transfer");
    cudaMemcpy(d_A, A, N * K * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, A, K * M * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_C, A, N * M * sizeof(float), cudaMemcpyHostToDevice);
    nvtxRangePop();

    nvtxRangePushA("Kernel Execution");
    // do operation
    nvtxRangePushA("Kernel A");
    for (int i = 0; i < n_iter; i++)
        sgemm_gpu_A(d_A, d_B, d_C, N, M, K, alpha, beta);
    cudaDeviceSynchronize();
    nvtxRangePop();    // Kernel A

    nvtxRangePushA("Kernel B");
    for (int i = 0; i < n_iter; i++)
        sgemm_gpu_B(d_A, d_B, d_C, N, M, K, alpha, beta);
    cudaDeviceSynchronize();

    nvtxRangePop();    // Kernel B
    nvtxRangePop();    // Kernel Execution
    cudaProfileStop();

在上面的代码中，我们已经扩大了关注的配置文件区域，以监视 NVTX 操作。我们还有Data Transfer、Kernel A、Kernel B和Kernel Execution作为 NVTX 范围。NVTX 支持多级注释，因此Kernel A和Kernel B范围将包含在Kernel Execution时间轴中。

要编译代码，我们应该为nvcc编译器提供-lnvToolsExt选项来提供 NVTX API 的定义。我们可以使用以下命令编译代码：

$ nvcc -m64 -gencode arch=compute_70,code=sm_70 -lnvToolsExt -o sgemm sgemm.cu

然后，NVIDIA 分析器可以在没有额外选项的情况下收集 NVTX 注释。我们可以使用以下命令对应用程序进行分析：

$ nvprof -f --profile-from-start off -o sgemm.nvvp ./sgemm.nvvp

以下屏幕截图显示了时间轴分析结果。在这个截图中，我们可以看到用绿色标记的标记和范围。这些绿色条有注释：

前面的屏幕截图为我们提供了以下信息：

我们可以通过 NVTX 注释来确定内存复制操作的位置。
我们可以通过包装区域来划分功能位置，例如kernel A和kernel B。
NVTX 注释可以堆叠多个级别的注释。正如我们所看到的，kernel A和kernel B包含在kernel execution注释中。

以下文件不仅介绍了 NVTX，还解释了如何使用 NVTX 来使用不同的颜色：devblogs.nvidia.com/cuda-pro-tip-generate-custom-application-profile-timelines-nvtx。NVTX 的一个应用是使用 NVTX 注释对深度学习网络进行分析。这提供了对网络操作瓶颈的洞察。我们将在本书的第十章《使用 CUDA 进行深度学习加速》中讨论这一点。

针对远程机器进行可视化分析

NVIDIA Visual Profiler 还可以分析远程应用程序。这个功能在远程应用程序开发时特别方便，尤其是在服务器端开发应用程序时。

有几种使用可视化分析器的方法，如下所示：

在主机上进行 CUDA 应用程序的分析
通过在目标端使用nvprof CLI 收集配置文件数据，将文件复制到主机并使用 Visual Profiler 打开
在目标平台上使用主机机器进行应用程序的分析

在主机机器上直接进行可视化分析非常方便，可以节省开发时间。此外，远程分析提供了与在主机机器上分析 GPU 应用程序相同的用户体验。唯一的例外是我们需要建立远程连接。主机管理的可视化分析提供的另一个好处是分析器会自动按需收集度量信息。

NVIDIA 分析器与主机机器中的 NVIDIA 分析器进行通信并收集分析数据。因此，您需要确认您的主机机器（台式机或笔记本电脑）应连接到远程机器。以下图显示了此连接的概述：

让我们尝试远程分析 GPU 应用程序。以下步骤介绍了如何在 NVIDIA Visual Profiler 中分析远程 GPU 应用程序：

首先，转到文件 | 新建会话。当您单击新建会话菜单时，您将看到以下对话框窗口：

然后，我们需要添加一个连接，方法是转到“管理连接...”菜单。然后，将出现“新的远程连接”对话框。通过单击“添加”按钮并在适当的部分输入远程机器信息来添加远程机器信息。然后，通过单击“完成”按钮关闭对话框。完成后，您将看到以下输出：

正如我们之前讨论的，主机和远程机器通过 SSH 进行通信，其默认端口号为 22。如果主机机器使用其他端口进行 SSH，您必须在新的远程会话创建对话框中通知它该端口号。

现在，我们需要通过单击 Toolkit/Script右侧的“管理...”按钮在远程机器上设置 CUDA Toolkit 路径。一个很好的开始是使用“检测”按钮。它会自动查找nvcc路径并自动设置配置信息。如果自动检测失败，您必须手动输入配置信息。完成配置过程后，单击“完成”按钮，如下所示：

通过单击“浏览”按钮在“文件”文本框的右侧指定 GPU 应用程序的二进制文件。它会要求您的远程机器登录密码。找到应用程序路径并设置应用程序路径。如果需要控制应用程序的行为，还可以输入应用程序的参数。完成应用程序和连接设置后，单击“下一步”按钮设置分析器的选项。
现在，我们将设置分析器选项。NVIDIA Visual Profiler 允许我们使用复选框设置分析器的选项，如下面的屏幕截图所示。单击“完成”，分析器将从应用程序收集分析数据：

您将在主机机器上看到时间线分析输出。

最后，分析分析时间线图的性能。单击要分析的任何内核函数。单击“执行内核分析”按钮；分析工具将收集相关的度量信息。通过这样做，您可以快速获得有关性能限制器的报告，并找到内核函数的瓶颈。

使用 CUDA 错误调试 CUDA 应用程序

具有专用的异常检查和检查错误是使软件具有高质量的基本特征之一。CUDA 函数通过返回每个函数调用的状态来报告错误。不仅如此，CUDA API，而且内核函数和 CUDA 库的 API 调用也遵循这个规则。因此，检测到重复错误是识别 CUDA 执行中错误的开始。例如，假设我们使用cudaMalloc()函数分配了全局内存，如下所示：

cudaMalloc((void**)&ptr, byte_size);

如果全局内存没有足够的空闲空间来分配新的内存空间怎么办？在这种情况下，cudaMalloc()函数返回一个错误来报告内存不足异常。通过使用cudaGetLastError()可以捕获由内核调用触发的标志。它返回记录的错误状态并重置标志的值。但是要小心处理这个标志：它的返回并不保证错误发生在 GPU 的最后执行，并且需要手动重置标志。

CUDA API 的返回值和cudaGetLastError()函数的返回值都是cudaError_t类型。这种cudaError_t类型是预定义的整数类型，应用程序可以识别发生了哪种类型的错误。例如，此类型定义如下：

Enum cudaErorr_t {
    cudaSuccess = 0,
    cudaErrorMemoryAllocation = 2, 
    cudaErrorUnknown = 30,
    cudaErrorNoDevice = 38,
    cudaErrorAssert = 59,
    cudaErrorTooManyPeers = 60,
    cudaErrorNotSupported = 71,
    ....
};

记住或翻译所有这些值是不切实际的。为此，CUDA 示例代码提供了一个辅助函数checkCudaError()，它位于common/inc/cuda_helper.h中。当 CUDA 函数返回错误时，此函数打印出错误消息。其函数定义如下：

#define checkCudaErrors(err) { \
    if (err != cudaSuccess) {  \

        fprintf(stderr, "checkCudaErrors() API error = %04d \"%s\" from file <%s>, line %i.\n", \
                err, cudaGetErrorString(err), __FILE__, __LINE__); \
        exit(-1); \
    } \
}
#endif

由于此函数被定义为宏，我们可以确定发生错误的行。

我们可以使用这个函数的两种方式。一种是在源代码中包含cuda_helper.h文件。另一种是将函数代码复制到代码中的某个位置。

然后，我们将使用checkCudaErrors()包装所有的 CUDA API 类，如下所示：

checkCudaErrors(cudaMalloc((void **)&d_A, N * K * sizeof(float)));
checkCudaErrors(cudaMalloc((void **)&d_B, K * M * sizeof(float)));
checkCudaErrors(cudaMalloc((void **)&d_C, N * M * sizeof(float)));

对于内核函数调用，我们将使用cudaGetLastError()函数来获取内核调用的错误标志，如下所示：

sgemm_kernel_A<<<dimGrid, dimBlock>>>(A, B, C, N, M, K, alpha, beta);
checkCudaErrors(cudaGetLastError());

然而，这段代码有一个问题：内核操作与主机异步，所以cudaGetLastError()只能捕获主机端的返回值。很可能错误是在应用程序的某个地方触发的。为了解决这种情况，您可以使用任何主机和设备同步函数；例如：

sgemm_kernel_A<<<dimGrid, dimBlock>>>(A, B, C, N, M, K, alpha, beta);
checkCudaErrors(cudaDeviceSynchronize());

现在，让我们通过修改源代码来测试错误检测代码。例如，您可以请求cudaMemcpy复制比分配大小更大的内存空间。在这种情况下，应用程序会返回一个错误消息，如下所示：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -lnvToolsExt -o sgemm ./sgemm.cu
CUDA error at sgemm.cu:93 code=11(cudaErrorInvalidValue) "cudaMemcpy(d_A, A, N * K * sizeof(float), cudaMemcpyHostToDevice)"

或者，您可以为 CUDA 内核传递一个NULL指针，以便内核访问无效的内存空间。在这种情况下，应用程序会在cudaDeviceSynchronize()中报告非法地址错误：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -lnvToolsExt -o sgemm ./sgemm.cu
CUDA error at sgemm.cu:104 code=77(cudaErrorIllegalAddress) "cudaDeviceSynchronize()"

这个错误检查宏非常有用，因为它报告了错误发生的源代码位置。然而，这个报告有一个缺点，就是它检测到的错误位置与实际发生的错误位置不匹配。

错误消息应该报告我们复制比分配的内存更大的内存位置导致非法值错误。因此，开发人员可以在内核调用之后立即识别错误消息。然而，这个错误检查代码只在主机上工作。因此，如果 GPU 操作没有正确同步，这可能会混淆 GPU 操作。例如，如果我们没有设置同步，只是检查错误，那么cudaDeviceSynchronize()函数可能会报告错误的位置。在这种情况下，我们可以设置CUDA_LAUNCH_BLOCKING=1环境变量，使所有内核执行与主机同步：

$ ./sgemm
CUDA error at sgemm.cu:104 code=77(cudaErrorIllegalAddress) "cudaDeviceSynchronize()" 
$ CUDA_LAUNCH_BLOCKING=1 ./sgemm
CUDA error at sgemm.cu:36 code=77(cudaErrorIllegalAddress) "cudaGetLastError()"

sgemm.cu的第 36 行是cudaGetLastError()调用，在sgemm内核调用之后。这就是我们放置一个预期错误的位置。我们可以在运行时确定正确的错误位置。

有两份官方文件可以帮助您了解不同类型的 CUDA 错误：

docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html
CUDA Toolkit 根路径中的include/driver_types.h

使用 CUDA assert 断言本地 GPU 值

即使您的 GPU 应用程序没有任何系统错误，您也需要检查计算结果，以确保执行的结果符合设计要求。为此，CUDA 提供了assert函数，它检查参数值是否为零。如果是，这个函数会引发一个错误标志，以便主机可以识别内核函数中存在错误。

断言用于验证操作结果是否符合预期。在 CUDA 编程中，可以从设备代码中调用assert函数，并在给定参数为零时停止内核的执行：

void assert(int expression);

这是assert函数的声明，与 C/C++的声明相同。当断言被触发时，应用程序会停止并报告其错误消息。如果应用程序由调试器启动，它会作为断点工作，以便开发人员可以调试给定的信息。例如，输出消息看起来像这样：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -lnvToolsExt -o sgemm ./sgemm.cu
sgemm.cu:29: void sgemm_kernel_A(const float *, const float *, float *, int, int, int, float, float): block: [16,64,0], thread: [0,0,0] Assertion `sum == 0.f` failed.

由于输出消息指向了确切的 CUDA 块和线程索引，开发人员可以轻松分析指定的 CUDA 线程的执行。

现在，让我们应用断言并看看它如何检测到预期的错误。我们将修改在GPU 应用程序中的性能优化目标范围部分中使用的 SGEMM 操作代码。

首先，在内核函数的中间放置断言代码。我们将看到表达式的效果，它应该是 false。断言代码可以编写如下：

__global__ void sgemm_kernel_A(const float *A, const float *B, float *C, int N, int M, int K, float alpha, float beta)
  {
      int col = blockIdx.x * blockDim.x + threadIdx.x;
      int row = blockIdx.y * blockDim.y + threadIdx.y;
      float sum = 0.f;
      for (int i = 0; i < K; ++i) 
          sum += A[row * K + i] * B[i * K + col];

      if (row == 0 && col == 0)
 assert(sum == 0.f);

      C[row * M + col] = alpha * sum + beta * C[row * M + col];
  }

您可以尝试其他索引值或尝试其他可能的错误。编译代码并运行它以查看输出。以下代码显示了此修改的输出错误：

sgemm.cu:29: void sgemm_kernel_A(const float *, const float *, float *, int, int, int, float, float): block: [0,0,0], thread: [0,0,0] Assertion `sum == 0.f` failed.

错误消息报告了断言触发的代码位置、内核函数的名称和 GPU 的线程索引。有了这些信息，我们可以很容易地找出应该从哪里开始分析。

实际上，assert函数的使用与普通 C/C++编程中的assert函数相同。一个区别是assert函数在设备代码中起作用。因此，它不仅报告事件位置和表达式，还显示块和线程索引。

然而，使用断言会对应用程序性能产生影响。因此，我们应该只在调试目的时使用断言。建议在生产环境中运行时禁用它。您可以通过在包含assert.h之前添加NDEBUG预处理宏来在编译时禁用断言。

使用 Nsight Visual Studio Edition 调试 CUDA 应用程序

对于 Windows 应用程序开发人员，CUDA Toolkit 提供了 Nsight Visual Studio Edition，它可以在 Visual Studio 中实现 GPU 计算。这个工具作为 Visual Studio 的扩展工作，但您可以构建、调试、分析和跟踪 GPU 应用程序以及主机。如果您的工作平台不是 Windows，则本节中的内容将不适用，您可以跳过它。

CUDA 调试器允许我们监视每个 CUDA 线程的 GPU 内核上的本地值。与普通主机调试一样，您可以在内核代码中设置断点并触发它们。您还可以设置条件，例如其他普通断点。有了这个功能，您可以为特定的 CUDA 线程索引触发断点并查看它们的本地变量。

这个工具可以与 CUDA Toolkit 一起安装。您可以从网站上获取最新版本。这不是强制性的，但是当您的开发环境使用旧的 CUDA Toolkit 和最新的 GPU 及其驱动程序时，建议使用它。访问 NVIDIA Nsight 网页（developer.nvidia.com/nsight-visual-studio-edition）下载并安装 Nsight。您需要 NVIDIA 开发人员会员资格才能获取该软件。您还需要安装推荐的显示驱动程序版本。

您可以通过转到 Visual Studio 菜单栏中的菜单 | Nsight 来找到 CUDA 工具。此菜单中有几个工具，其中一些如下：

图形调试：用于图形（Direct3D、OpenGL 和 Vulkan）应用程序的调试器
CUDA 调试（Next-Gen）：用于同时调试 CPU 和 GPU 代码的调试器（Turing、Volta 和 Pascal 与最新驱动程序）
CUDA 调试（传统）：仅用于 GPU 内核的调试器（具有旧驱动程序的 Pascal、Maxwell 和 Kepler）
性能分析：用于分析当前 GPU 应用程序的性能
CUDA 内存检查器：用于在运行时检查 GPU 内存违规（如前一节中介绍的）

在本节中，我们将重点放在 CUDA 调试（Next-Gen）上。这是因为 Next-Gen 调试器可以支持包括 Turing 和 Volta 在内的最新架构。CUDA 内存检查器将在本章末尾介绍。

现在，让我们配置一个示例项目，并看看我们如何使用 Nsight Visual Studio Edition 调试应用程序。您可以使用默认的示例代码，或者用我们之前介绍的 CUDA 代码替换代码。您还可以使用05_debug/05_debug_with_vs文件中提供的示例代码。这是一些简单的 SAXPY 代码。

将项目属性设置为生成适当的设备目标代码。在项目的属性页面中，您可以指定目标代码版本。在 CUDA C/C++ | 代码生成文本框中列出您想要在其中使用的架构版本：

上述截图显示了 CUDA 设备代码生成属性页面。您可以设置几个nvcc选项，例如目标 GPU 的计算能力、每个线程的寄存器限制以及在编译时冗长的 CUDA 内核信息。

在第 34 行和第 75 行设置断点，其中第 34 行是内核函数的中间位置，第 75 行是从主机复制数据到设备的位置。然后，使用以下方法之一编译并开始调试：

在 Visual Studio 菜单栏中导航到 Nsight，然后单击“开始 CUDA 调试（Next-Gen）”。
在“解决方案资源管理器”中右键单击项目，选择“调试|开始 CUDA 调试（Next-Gen）”。
转到 Nsight CUDA 调试工具栏，单击“开始 CUDA 调试（Next-Gen）”。

Windows 防火墙可能会询问您是否信任并允许 Nsight 的网络连接。这是正常的，因为 Nsight 使用内部网络来监视 GPU 设备。单击“接受”并继续调试。当前的 Nsight Visual Studio Edition 提供了两种调试选项。这取决于目标 GPU 架构版本。如果您的 GPU 是 Volta 或 Turing，建议使用“Next-Gen”调试。如果您的 GPU 是 Pascal，则适当的调试器取决于驱动程序版本。为了澄清，请访问 NVIDIA 支持的 GPU 列表：developer.nvidia.com/nsight-visual-studio-edition-supported-gpus-full-list。

应用程序将在应用程序启动的地方停止。继续跟踪。应用程序将在主机的第 75 行和设备的第 34 行停止。从中我们可以了解到，Nsight 可以同时跟踪主机和设备上的 GPU 应用程序。

当黄色箭头停在内核函数中时，您可以查看局部变量。全局索引中的线程索引为0。由于 CUDA 并行发出多个 CUDA warp 和 CUDA 线程，因此您可以通过更改blockIdx和threadIdx来查看其他线程的局部变量。基本的 CUDA 线程调试控制单元是 warp。换句话说，您可以控制调试器以遍历活动 warp。Nsight 调试器在 Nsight 菜单栏中的“上一个活动 warp/下一个活动 warp”菜单中提供了此功能。

以下屏幕截图显示了我们在调试时出现的 Nsight 调试控件：

如果更改 warp，您会发现在“Autos”面板中监视的局部变量会随着 warp 的变化而更新索引。例如，以下屏幕截图显示了“Autos”窗口，该窗口报告了活动 warp 中所选线程的局部变量，即正在由主导线程监视的局部变量的值：

选择的线程更改后，Autos 值会更新。以下屏幕截图显示了通过移动到下一个活动 warp 所做的更改：

Next-Gen CUDA 调试器提供了三种类型的窗口——warp info、lanes 和 GPU registers。黄色箭头表示当前的 GPU 执行，并以三个方面显示其信息：

Warp Info 窗口提供了另一种选择活动 warp 的方法。您可以在菜单栏中从 Nsight | Window | Warp Info 打开该窗口。窗口如下所示：

每行表示 CUDA 网格中的活动 warp。第四列“Shader Info”显示了每个 warp 的块和主导线程索引。第五列“threads”显示了 warp 中 CUDA 线程的状态。单元格的颜色表示每个线程的状态。由于我们在断点处观察它们，它们都是红色的，但在调试过程中您会看到其他颜色。以下屏幕截图解释了每种颜色在线程状态方面的含义：

双击任何 warp，查看 autos 窗口中的局部变量是如何更新的。

Lanes 窗口允许您在所选活动 warp 内选择特定的 CUDA 线程。一个 lane 指的是 warp 中的一个线程。您可以从 Nsight | Window | Lanes 中打开该窗口。通过双击一个 lane，您可以发现 autos 窗口中的局部变量根据更新的索引而更新：

活动 warp 中的 lanes 窗口信息。

寄存器窗口显示了 GPU 寄存器的当前状态。如果它们的值被更新，它们将变为红色。

如果您想了解如何使用 Nsight Visual Studio Edition，请阅读 NVIDIA 官方用户指南。它介绍了如何配置调试环境，如何使用它，以及各种情况下的详细提示。 (docs.nvidia.com/nsight-visual-studio-edition/Nsight_Visual_Studio_Edition_User_Guide.htm)。

使用 Nsight Eclipse Edition 调试 CUDA 应用程序

对于 Linux 和 OSX 平台开发，CUDA Toolkit 提供了 Nsight Eclipse Edition。这个工具基于 Eclipse，因此开发人员可以很容易地在 CUDA C 开发中使用这个工具。

Nsight Eclipse Edition 是基于 Eclipse 用于 CUDA 应用程序开发的。您可以使用它来编辑、构建、调试和分析 CUDA 应用程序。它使得在 Linux 和 OSX 中进行 CUDA C/C++开发变得简单。这个工具作为 CUDA Toolkit 的一部分安装，因此您不必单独安装这个工具。但是，如果您使用 Linux，需要配置 Java 7 才能使用它。

Nsight Eclipse Edition 是基于 Eclipse 4.4.0 版本（2014 年发布的 Luna 版本）构建的，并且基于 Java 7 构建。

Nsight 可以通过终端中的nsight命令或者 X 窗口应用程序列表中执行。

现在，让我们从终端或 X 窗口桌面打开 Nsight，以便我们可以编译和分析给定的示例。要么创建一个新的 CUDA 项目，要么打开05_debug/06_debug_with_eclipse中提供的示例项目。如果要创建项目，请选择 CUDA C/C++项目。空项目只会给您一个空项目，而 CUDA Runtime 项目会给您一个带有一些示例代码的项目。如果要使用示例项目，请使用文件 | 导入 | 导入现有项目到工作区。

现在，让我们在sgemm内核函数中设置一个断点。就像在 Eclipse 中的普通 C/C++项目一样，您可以在nsight中构建和调试 CUDA 应用程序。在内核函数的起始点（第 23 行）设置一个断点，如下所示：

对于内核函数调试来说，一个很好的起点是在线程索引计算之后。设置一个断点来暂停 GPU 的执行。现在，通过单击菜单面板中的绿色 bug 来编译和开始调试。在调试窗口切换调试透视之时，点击继续，直到达到我们设置的断点。

Nsight 允许您监视活动 warp 中的局部变量和寄存器。首先，它会在 CUDA 网格中的领先 CUDA 线程（CUDA 线程0）处停止应用程序。然后，您可以从调试窗口切换到其他 CUDA 活动 warp，并使用 CUDA 窗口检查每个 CUDA 线程，就像这样：

以下截图显示了所选 CUDA 线程的局部变量信息。Nsight 会在这些值更新时更新它们：

上述截图显示了 Eclipse 的调试透视窗口中的 Debug 窗口和 CUDA 窗口。调试窗口提供了在所选 GPU 上的活动 warp 中进行 CUDA warp 选择的功能，并且可以在所选活动 warp 内进行 lane 选择。

NVIDIA 还有一个 Nsight Eclipse Edition 用户指南。您可以通过访问docs.nvidia.com/cuda/nsight-eclipse-edition-getting-started-guide/index.html来了解更多关于这个工具的信息。

使用 CUDA-GDB 调试 CUDA 应用程序

CUDA 工具包提供了 CUDA-GDB，它支持 CUDA C/C++调试，用于诸如 C/C++ GDB 之类的程序。这对于直接调试没有 X 窗口环境或远程调试的 CUDA C/C++应用程序非常有用。

要调试 GPU 应用程序，Makefile应该包括主机的-g调试标志和 GPU 的-G调试标志。基本上，CUDA 的 GDB 用法与主机调试相同，只是在 CUDA 操作之外还有一些额外的调试功能。例如，我们可以设置特定的 CUDA 线程和 CUDA 感知断点。

CUDA-GDB 的断点

让我们看看cuda-gdb如何帮助我们检测代码中的错误。我们将在代码中设置断点，并查看主机和 GPU 上的局部值。为此，将工作目录切换到05_debug/07_debug_with_gdb目录。我们将通过将其与适当的行匹配来检查cuda-gdb的操作。

首先，让我们使用以下命令编译源代码：

$ nvcc -run -m64 -g -G -Xcompiler -rdynamic -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o simple_sgemm ./simple_sgemm.cu

然后，我们应该执行cuda-gdb，这样我们就可以在终端上调试应用程序，如下所示：

$ cuda-gdb simple_sgemm

我们可以在代码的特定行上设置断点，如下所示：

(cuda-gdb) break simple_gemm.cu:21

或者，我们可以按照内核函数的名称设置断点，如下所示。这将在函数的入口点触发断点：

(cuda-gdb) break sgemm_kernel

如果cuda-gdb警告指出断点希望在未来的共享库加载时挂起，则回答y。您也可以在主机代码上设置断点。

使用断点的一个问题是，断点将根据 CUDA 线程的数量触发。因此，我们应该提供条件信息，以便针对特定的 CUDA 线程设置断点。条件断点如下：

(cuda-gdb) break sgemm_kernel if blockIdx.y == 2

当然，我们可以修改预定义断点的条件如下：

(cuda-gdb) cond 3 // break 3 is defined previously

让我们使用run命令执行示例应用程序。如果应用程序遇到任何断点，CUDA-GDB 将提供有关它的信息。以下代码显示了应用程序在第21行遇到断点时cuda-gdb的报告：

(cuda-gdb) run
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (5,0,0), device 0, sm 0, warp 0, lane 5]
Thread 1 "simple_sgemm" hit Breakpoint 1, sgemm_kernel<<<(128,128,1),(16,16,1)>>> (A=0x7fffb6000000, B=0x7fffb7000000, C=0x7fffb4000000, N=2048, M=2048, K=2048, alpha=2, beta=1) at simple_sgemm.cu:21
21 int col = blockIdx.x * blockDim.x + threadIdx.x;

现在，是时候使用 GDB 命令来跟踪代码或监视活动变量了。我们可以使用 next（或n）、step（或s）、continue（或c）和 finish（或fin）来跟踪内核函数。然而，当我们到达内核代码的末尾并需要在主机和设备之间切换目标硬件时，我们应该使用continue命令。

使用 CUDA-GDB 检查变量

除了默认的 GDB 命令之外，CUDA-GDB 提供了可以与 CUDA 内核一起使用的调试功能。以下是您可以使用 CUDA-GDB 做的事情。

列出内核函数

与普通函数一样，CUDA-GDB 可以在内核函数上设置断点。一旦应用程序被断点停止，您可以列出它们如下：

(cuda-gdb) info cuda kernels
Kernel Parent Dev Grid Status   SMs Mask     GridDim  BlockDim Invocation
*      0      -   0    1 Active 0xffffffff (128,128,1) (16,16,1) sgemm_kernel(A=0x7ffff5a79010, B=0x7ffff4a78010, C=0x7ffff3a77010, N=2048, M=2048, K=2048, alpha=2, beta=1)

正如您所看到的，前面的输出显示了内核的配置信息和输入参数变量。

变量调查

CUDA-GDB 帮助我们通过选择特定的线程块索引和线程索引来跟踪特定的 CUDA 线程。有了这个功能，您可以将当前焦点移动到指定的线程。在这个例子中，块大小为 16，col变量被定义为x维度上的 CUDA 线程索引。以下代码显示了 CUDA-GDB 如何通过更改线程索引来报告所选的局部变量的值：

(cuda-gdb) print col
$1 = <optimized out>
(cuda-gdb) cuda kernel 0 block 1,2,0 thread 3,4,0
21 int col = blockIdx.x * blockDim.x + threadIdx.x;
(cuda-gdb) s
22 int row = blockIdx.y * blockDim.y + threadIdx.y;
(cuda-gdb) p col
$2 = 19

检查当前焦点线程的信息：

(cuda-gdb) cuda device kernel block thread
kernel 3, block (1,2,0), thread (3,4,0), device 0

有了手头的信息，我们可以追踪 CUDA 线程。

如果您想了解有关 CUDA-GDB 的更多信息，请查看 NVIDIA 的用户指南文档：docs.nvidia.com/cuda/cuda-gdb/index.html。

使用 CUDA-memcheck 进行运行时验证

CUDA 编程的一个困难点是处理内存空间。由于 CUDA 线程并行操作，边界条件或意外的索引操作可能会违反有效的内存空间。CUDA memcheck 是一个运行时测试工具，如果任何 GPU 操作超出了无效的内存空间，它将验证内存访问。该工具检测以下内存错误：

名称	位置	描述	精确
内存访问错误	设备	无效的内存访问（超出边界，未对齐）	O
硬件异常	设备	硬件错误	X
Malloc/free 错误	设备	在 CUDA 内核中不正确使用`malloc()`/`free()`	O
CUDA API 错误	主机	CUDA API 的错误返回	O
cudaMalloc 内存泄漏	主机	使用`cudaMalloc()`分配的设备内存未被应用程序释放	O
设备堆内存泄漏	设备	在设备代码中使用`malloc()`分配的设备内存未被应用程序释放	X

精确（O）表示 memcheck 可以指定崩溃的行和文件。另一方面，不精确（X）表示该工具可以识别错误，但由于并发状态，无法指定错误点。cuda-memcheck不需要重新编译进行测试。但是，如果我们使用一些额外的nvcc选项进行编译，我们可以跟踪错误点。nvcc选项包括生成行号信息的-lineinfo和用于保留函数符号的-Xcompiler -rdynamic。

基本上，cuda-memcheck是一个独立的工具，可以在运行时验证 GPU 应用程序。以下命令显示了它在独立模式下的格式：

$ cuda-memcheck [options] <application>

这个工具也可以与 CUDA-GDB 一起使用，帮助开发人员识别错误并进行调试。在 CUDA-GDB 命令行中，使用set cuda memcheck on命令启用内存检查。这样，CUDA-GDB 可以识别与内存相关的异常。

检测内存越界

现在，让我们看看cuda-memcheck如何检测内存异常并与 CUDA-GDB 一起工作。为了简化这个过程，我们将编写一些错误的代码，并查看cuda-memcheck如何报告结果。让我们从一些干净的代码开始。您可以使用05_debug/08_cuda_memcheck中提供的示例代码进行测试。让我们使用cuda-memcheck测试代码并验证它：

$ nvcc -m64 -g -G -Xcompiler -rdynamic -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o simple_sgemm ./simple_sgemm.cu
$ cuda-memcheck simple_sgemm
========= CUDA-MEMCHECK
Application finished successfully.========= ERROR SUMMARY: 0 errors

现在，让我们将一些错误的代码放入内核函数中，如下所示。如果您愿意，您也可以放入其他错误：

For instance, you may add one to the row value.
__global__ void sgemm_kernel(const float *A, const float *B, float *C, int N, int M, int K, float alpha, float beta)
{
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    row += 1;

    float sum = 0.f;
    for (int i = 0; i < K; ++i)
        sum += A[row * K + i] * B[i * K + col];
    C[row * M + col] = alpha * sum + beta * C[row * M + col];
}

让我们编译并启动代码。内核将返回一个 CUDA 错误，checkCudaErrors()将报告一个错误消息，如下所示：

CUDA error at simple_sgemm_oob.cu:78 code=77(cudaErrorIllegalAddress) "cudaDeviceSynchronize()"

然而，如果我们希望确定内核代码中的哪一行是问题的根本原因，这些信息是不够的。使用cuda-memcheck，我们可以确定哪个 CUDA 线程和内存空间触发了错误，并给出堆栈地址：

$ cuda-memcheck simple_sgemm_oob

输出如下：

前面的屏幕截图显示了cuda-memcheck独立执行的一部分，显示了内核中检测到的所有错误。在这种情况下，cuda-memcheck报告检测到在第 27 行发生的内存违规错误。默认情况下，cuda-memcheck在检测到错误时会停止应用程序的执行。

在这种情况下，我们可以通过检查相关变量来轻松找到根本原因，使用cuda-gdb。为此，我们需要使用cuda-gdb启动应用程序，并启用cuda-memcheck，如下所示：

$ cuda-gdb simple_sgemm_oob
(cuda-gdb) set cuda memcheck on
(cuda-gdb) run

这个过程使cuda-gdb从cuda-memcheck报告非法内存访问检测：

前面的屏幕截图显示了cuda-gdb与cuda-memcheck的报告。开发人员可以轻松地确定simple_sgemm_oob.cu中的第 27 行触发了报告的错误。根据给定的信息，我们可以开始调查哪一块内存访问了无效的空间，如下所示：

(cuda-gdb) print A[row * K + i]
Error: Failed to read generic memory at address 0x7fffc7600000 on device 0 sm 41 warp 20 lane 16, error=CUDBG_ERROR_INVALID_MEMORY_SEGMENT(0x7).
(cuda-gdb) print row * K + i
$1 = 4194304

在不费吹灰之力的情况下，我们可以确定访问A[row * K + i]会触发错误，并且请求的值超出了全局内存（A）的分配空间。通过这种方式，您可以轻松地缩小根本原因。

检测其他内存错误

CUDA memcheck 工具提供了额外的软件验证功能，其中一些如下：

名称	描述	选项
内存泄漏	用于识别内存泄漏	`--leak-check full`
竞争检查	用于分析多个线程之间对共享内存的冲突访问的竞争危险	`--tool racecheck`
初始化检查	在没有初始化的情况下识别设备全局内存访问	`--tool initcheck`
同步检查	验证同步原语的正确使用，如`__syncthreads()`，`__syncwarp()`和协作组 API	`--tool synccheck`

这些工具假设内存访问是正确的或经过验证的，并且不检查内存错误。因此，您需要确认您的应用程序中不存在内存错误。其他有用的 memcheck 选项包括--save，我们可以用它来将输出保存到磁盘，以及--print-level，我们可以用它来控制输出的详细级别。

NVIDIA 为cuda-memcheck提供了用户指南。该文档将帮助您使用 GPU 验证您的应用程序并检测意外错误 (docs.nvidia.com/cuda/cuda-memcheck/index.html)。

使用 Nsight Systems 对 GPU 应用程序进行分析

在本节中，我们将介绍新引入的 CUDA 分析器工具，即 Nsys 和 Nvprof。这些分析器支持 Volta 架构及更高版本的 GPU。它是图灵架构 GPU 中的主要分析器。我们将先介绍 Nsys，然后在下一节介绍 Nvprof。

Nsight Systems (developer.nvidia.com/nsight-systems)是一个系统范围的性能分析工具，可以在时间轴上可视化操作并轻松找到优化点。在时间轴分析方面，Nsight Systems 提供了系统利用率信息，以便我们可以分析瓶颈点。我们可以从 NVIDIA 网站获取 Nsight Systems，但 CUDA 10 默认包含了 Nsight Systems 在工具包中。我们只需要确保它安装正确即可。

对于 CLI，我们应该设置PATH以便于我们的操作，因为它的路径与普通的 CUDA 二进制文件分开。我们可以使用以下命令将其包含在PATH环境变量中：

export PATH=$PATH:/usr/local/cuda/bin:/usr/local/cuda-10.1/NsightSystems-2019.3/Target-x86_64/x86_64

Nsys 提供了两个接口：一个用于 GUI，一个用于 CLI。在主机上，我们可以通过 GUI 运行应用程序来收集应用程序的采样信息。在远程机器上，我们可以通过 CLI 收集分析数据，使用以下命令：

$ nsys profile -t osrt,cuda,nvtx,cublas,cudnn -o baseline -w true <command>

这个选项可以解释如下：

	选项	开关
跟踪	`-t`/`--trace`	`cuda`: 用于跟踪 CUDA 操作，`nvtx`: 用于跟踪`nvtx`标签，`cublas`, `cudnn`, `opengl`,`openacc`: 用于跟踪 API 操作，`osrt`: 用于跟踪 OS 运行时库，`none`: 不进行 API 跟踪
输出文件	`-o`/`--output`	输出文件名
显示输出	`-w`/`--show-`output	`true`/`false`: 在终端上打印出分析器的行为

例如，我们可以从02_nvtx SGEMM 应用程序中获得一个名为sgemm.qdrep的分析文件。让我们比较 Nsight Systems 和 NVIDIA Visual Profiler 之间的分析输出。我们可以使用以下命令收集 Nsys 的分析数据：

$ nsys profile -t osrt,cuda,nvtx -o sgemm -w true ./sgemm

这是来自 Nsys 的分析时间轴视图：

以下屏幕截图显示了来自 NVIDIA Visual Profiler 的分析时间轴视图：

Visual Profiler 显示操作事件块，而 Nsight Systems 同时显示系统利用率。因此，我们可以轻松地看到哪些资源（CPU 核心、GPU 或 PCIe 总线）对性能产生影响。此外，Nsight Systems 提供了更具交互性的性能分析体验。当双击任何函数操作时，Nsight Systems Viewer 会展开时间轴以适应窗口，并帮助我们检查操作。此外，Nsight Systems 使我们能够轻松地发现在某个 NVTX 区域下发生的内核执行次数。在 Visual Profiler 时间轴视图中，内核执行看起来像是单个执行，但 Nsight Systems 显示了分离的执行。

现在我们已经确定了应该优化的函数，我们可以继续使用 Nsight Compute，这是另一个新的性能分析器，用于检查内核函数的 GPU 操作。

使用 Nsight Compute 进行内核性能分析

Nsight Compute 是一个用于计算的内核级性能分析器。它收集 GPU 指标信息，并帮助我们专注于 CUDA 内核的优化。换句话说，这个工具涵盖了 Visual Profiler 的性能分析功能。

Nsight Compute 提供两种接口：GUI 和 CLI。GUI 支持主机和远程应用程序性能分析，而 CLI 适用于目标机器。然而，我们可以使用 GUI 获取分析数据并查看结果。

使用 CLI 进行性能分析

为了方便使用 Nsight Compute CLI，我们需要在/usr/local/cuda-10.1/NsightCompute-2019.3/nv-nsight-cu-cli中设置PATH环境变量。然后，我们可以使用以下命令收集性能分析数据：

$ nv-nsight-cu-cli -o <output filename> <application command>

这个命令收集 GPU 执行指标信息，并将数据保存到指定的文件中。如果我们没有提供输出文件名，Nsight Compute 将把收集到的指标报告输出到控制台，从而在控制台上提供快速的指标性能报告。

由于我们可以指定性能分析目标，我们可以限制 Nsight Compute 收集以下信息：

--kernel-regex：指定要进行性能分析的内核
--设备：专注于对特定 GPU 进行性能分析

当我们需要在控制台上查看报告时，这个功能非常有用。

使用 GUI 进行性能分析

通过在 Nsight Compute 中打开一个新项目，我们可以启动性能分析操作。以下截图显示了性能分析配置。对于主机应用程序开发，请连接到本地主机。或者，您可以指定要进行性能分析的目标 GPU 服务器：

当然，我们也可以打开使用 CLI 工具在目标机器上生成的nsight-cuprof-report文件。例如，我们可以使用以下命令创建 sgemm 性能分析文件：

$ nv-nsight-cu-cli -o reduction reduction

对于 OSX 用户，Nsight Systems 将需要目标glib库进行远程性能分析。在这种情况下，我们应该从 Nsight Compute 安装映像中复制该库。它将所需的库提供为一个名为 target 的目录，并将该目录复制到Applications/NVIDIA Nsight Compute.app/target目录。

为了方便起见，我们将使用来自第三章 CUDA 线程编程的减少示例代码。它有两个不同寻址的并行减少实现。您可以在03_cuda_thread_programming/05_warp_divergence目录中找到代码。完成连接和应用程序可执行文本栏的设置后，单击启动按钮。然后，按下Ctrl + I，Ctrl + K键以运行到下一个内核函数，然后性能分析器将停在reduction_kernel_1处。按下Ctrl + I，Ctrl + P键以对此内核进行性能分析。然后您将得到以下输出。这张图片展示了 Nsight Compute 基于 GUI 的第一个内核性能分析：

显示基于 GUI 的配置文件（用于第一个内核配置文件）

它提供了交互式配置文件和调试。使用步骤控制调试按钮，我们可以调试 CUDA API 和内核函数。我们还可以使用左侧 API 流面板上的控制按钮移动到下一个内核函数或下一个配置文件范围。在右侧面板上，您可以获取内核的详细配置文件信息。

我们还可以通过启用自动配置文件来自动获取配置文件结果，具体操作如下：转到菜单栏，选择 Profile | Auto Profile。然后，继续进行应用程序。Nsight Systems 将配置所有的内核函数。或者，您可以通过单击窗口顶部的 Profile Kernel 按钮来手动配置内核函数。当我们使用 CLI 收集的配置文件结果时，我们将只看到所有内核函数的配置文件数据。

性能分析报告

在交互式配置文件窗口的右侧面板上，我们可以看到 Nsight Compute 提供了性能分析报告。从报告中，我们可以确定性能限制因素并调查未充分利用的资源。此外，Nsight Compute 还根据资源利用统计数据提供优化建议。我们也可以直接从直接配置文件中识别它们。

此外，Nsight Compute 通过分析 GPU 组件的利用率提供优化建议。它找到瓶颈并建议进行推荐的调查以优化内核。

此报告页面提供了每个组件的利用率，如计算、内存、调度器、指令、warp 等。此外，您可以通过扩展每个组件的左上箭头来获取更多详细信息。以下图片显示了内存工作负载分析的示例报告：

在 Nsight Compute 中，我们可以轻松获取这样的详细信息。在以前的分析器 NVIDIA Profiler 中，我们应该执行每个分析以获取这样的信息。

基线比较

在优化过程中，我们应该将新结果与基线操作进行比较。为了使这项任务对我们来说更容易，Nsight Compute 提供了基线比较功能。单击性能报告面板顶部的 Add baseline 按钮，并将其更改为其他内核函数。然后，我们可以使用 Nsight Compute 来比较内核函数的利用率。以下屏幕显示了这一点：

内核函数利用率的比较

如果我们希望追踪我们的优化工作并确定有效的组件，这将非常有用。

源视图

Nsight Compute 提供了各种我们可以调查的页面。其中一个有用的页面是 Source 页面。如果 CUDA 应用程序是使用-lineinfo选项构建的，Nsight Compute 可以显示与 CUDA C/C++源代码相关的信息和 CUDA SASS 代码。然后，我们可以分析瓶颈代码并调查它与 SASS 代码级别的关系。此外，它提供了一个 Live Registers 数字，以便我们可以调查内核函数中所需寄存器的数量。以下截图显示了 Source 页面：

如果您需要了解更多关于此功能的信息，您可以在此文档中找到相关信息-docs.nvidia.com/nsight-compute/NsightCompute/index.html#profiler-report-source-page。

Nsight Compute 提供了一个以 CUDA 内核性能分析为中心的操作，我们可以用来验证 Night Systems 和 Nsight Compute 具有不同的优化范围。

总结

在本章中，我们已经介绍了如何配置 GPU 应用程序并对其进行调试。了解这些 CUDA 工具将有助于您高效和有效地开发，因为它们可以帮助您找到瓶颈，并在短时间内找到错误和漏洞。

到目前为止，我们一直专注于单个 GPU 应用程序开发。然而，许多 GPU 应用程序使用多个 GPU 来实现更好的性能。在下一章中，我们将介绍如何编写能在多个 GPU 上运行并且具有可扩展性性能的代码。您将学习什么因素会影响性能以及如何实现良好的性能水平。您还将能够应用本章涵盖的工具来加强多 GPU 系统及其经验，解决下一章的问题。

第六章：可扩展的多 GPU 编程

到目前为止，我们一直致力于在单个 GPU 上获得最佳性能。密集节点与多个 GPU 已成为即将到来的超级计算机的迫切需求，特别是自从 ExaFLOP（每秒千亿次操作）系统成为现实以来。 GPU 架构具有高能效，因此近年来，具有 GPU 的系统在 Green500 榜单（www.top500.org/green500）中占据了大多数前十名。在 2018 年 11 月的 Green500 榜单中，前十名中有七个基于 NVIDIA GPU。

NVIDIA 的 DGX 系统现在在一个服务器中有 16 个 V100 32GB。借助统一内存和诸如 NVLink 和 NvSwitch 之类的互连技术，开发人员可以将所有 GPU 视为一个具有 512GB 内存的大型 GPU（16 个 GPU *每个 32GB）。在本章中，我们将深入讨论编写 CUDA 代码的细节，并利用 CUDA-aware 库在多 GPU 环境中实现节点内和节点间的可伸缩性。

在本章中，我们将涵盖以下主题：

使用高斯消元法解线性方程
GPUDirect 点对点
MPI 简介
GPUDirect RDMA
CUDA 流
额外的技巧

技术要求

本章需要一台带有现代 NVIDIA GPU（Pascal 架构或更高版本）的 Linux PC，并安装了所有必要的 GPU 驱动程序和 CUDA Toolkit（10.0 或更高版本）。如果您不确定您的 GPU 架构，请访问 NVIDIA GPU 网站（developer.nvidia.com/cuda-gpus）并确认您的 GPU 架构。本章的代码也可以在 GitHub 上找到：github.com/PacktPublishing/Learn-CUDA-Programming。

本章中的示例代码是使用 CUDA 版本 10.1 开发和测试的。但是，建议您使用最新版本（CUDA）或更高版本。

由于本章需要展示多 GPU 的交互，我们需要至少两个相同类型和架构的 GPU。还要注意，一些功能，如 GPUDirect RDMA 和 NVLink，仅支持 NVIDIA 的 Tesla 卡。如果您没有像 Tesla P100 或 Tesla V100 这样的 Tesla 卡，不要灰心。您可以安全地忽略其中一些功能。与我们在这里展示的情况相比，性能数字将会有所变化，但相同的代码将仍然有效。

在下一节中，我们将看一个示例，使用流行的高斯算法解决一系列线性方程，以演示如何编写多 GPU。

使用高斯消元法解线性方程

为了演示在节点内和节点间使用多个 GPU，我们将从一些顺序代码开始，然后将其转换为节点内和节点间的多个 GPU。我们将解决一个包含M个方程和N个未知数的线性方程组。该方程可以表示如下：

A × x = b

在这里，A是一个具有M行和N列的矩阵，x是一个列向量（也称为解向量），具有N行，b也是一个具有M行的列向量。找到解向量涉及在给定A和b时计算向量x。解线性方程组的标准方法之一是高斯消元法。在高斯消元法中，首先通过执行初等行变换将矩阵A减少为上三角矩阵或下三角矩阵。然后，通过使用回代步骤解决得到的三角形方程组。

以下伪代码解释了解线性方程所涉及的步骤：

1\. For iteration 1 to N (N: number of unknowns) 
    1.1 Find a row with non-zero pivot
    1.2 Extract the pivot row
    1.3 Reduce other rows using pivot row
2 Computing the solution vector through back substitution

让我们看一个示例，以便理解算法。假设方程组如下：

首先，我们将尝试设置基线系统，如下所示：

准备您的 GPU 应用程序。此代码可以在本书的 GitHub 存储库中的06_multigpu/gaussian文件夹中找到。
使用nvcc编译器编译您的应用程序，如下所示：

$ nvcc -o gaussian_sequential.out gaussian_sequential.cu
$ nvcc -o gaussian_single_gpu.out gaussian_single_gpu.cu
$ $ time ./gaussian_sequential.out
$ time ./gaussian_single_gpu.out

前面的步骤编译并运行了本章中存在的两个版本的代码：

顺序运行的 CPU 代码
在单个 GPU 上运行的 CUDA 代码

现在，让我们看看高斯消元的单 GPU 实现中的热点。

高斯消元的单 GPU 热点分析

让我们尝试理解和分析顺序和单 GPU 代码以建立基线。在此基础上，我们将增强并添加对多 GPU 运行的支持。

顺序 CPU 代码：以下代码显示了顺序实现的提取代码：

for( int n = 0; n < N; n++ ){
// M: number of equations, N: number of unknowns
    for( int pr = 0; pr < M; pr++ ){
        // finding the pivot row 
        //if pr satisfies condition for pivot i.e. is non zero 
        break; 
    }
    for( int r = 0; r < M; r++ ){
        // reduce all other eligible rows using the pivot row
        double ratio = AB[r*N+n]/AB[pr*N+n]
        for( int nn = n; nn < N + 1; nn++ ){
            AB[r * N + nn] -= (ratio*AB[pr * N + nn]);
        }
    }
}

从视觉上看，发生的操作如下：

在这里，高斯消元中的行数等于方程的数量，列数等于未知数的数量。在前面的图表中显示的pr行是主元行，将用于使用主元素减少其他行。

我们可以做出的第一个观察是，我们正在对增广矩阵进行操作，将A矩阵与b向量合并。因此，未知数的大小为N+1，因为增广矩阵的最后一列是b向量。创建增广矩阵有助于我们只处理一个数据结构，即矩阵。您可以使用以下命令对此代码进行分析。分析结果将显示guassian_elimination_cpu()函数完成所需的时间最长：

$ nvprof --cpu-profiling on ./guassian_sequential.out

CUDA 单 GPU 代码：通过前几章的学习，我们期望您已经熟悉了如何编写最佳的 GPU 代码，因此我们不会详细介绍单个 GPU 实现。以下摘录显示，在单个 GPU 实现中，三个步骤被称为三个用于找到N未知数的核心：

findPivotRowAndMultipliers<<<...>>>：该核心查找主元行和乘数，应用于行消除。
extractPivotRow<<<>>>：该核心提取主元行，然后用于执行行消除。
rowElimination<<<>>>：这是最终的核心调用，在 GPU 上并行进行行消除。

以下代码片段显示了数据在复制到 GPU 后迭代调用的三个核心：

<Copy input augmented matrix AB to GPU>
...
for( int n = 0; n < N; n++ ){
// M: number of equations, N: number of unknowns
    findPivotRowAndMultipliers<<<...>>>(); 
    extractPivotRow<<<...>>>(); 
    rowElimination<<<...>>>(); 

}

本章的重点是如何增强此单个 GPU 实现以支持多个 GPU。但是，为了填补 GPU 实现中的缺失部分，我们需要对单个 GPU 实现进行一些优化更改：

高斯消元算法的性能受内存访问模式的影响很大。基本上，它取决于 AB 矩阵的存储方式：
找到主元行更喜欢列主格式，因为如果矩阵以列主格式存储，则提供了合并访问。
另一方面，提取主元行更喜欢行主格式。
无论我们如何存储AB矩阵，内存访问中都无法避免一个合并和一个跨步/非合并的访问。
列主格式对于行消除核心也是有益的，因此对于我们的高斯消元核心，我们决定存储 AB 矩阵的转置而不是 AB。AB 矩阵在代码开始时通过transposeMatrixAB()函数转置一次。

在下一节中，我们将启用多 GPU P2P 访问并将工作分配给多个 GPU。

GPU 直接点对点

GPUDirect 技术是为了允许 GPU 在节点内部和跨不同节点之间进行高带宽、低延迟的通信而创建的。该技术旨在消除一个 GPU 需要与另一个 GPU 通信时的 CPU 开销。GPUDirect 可以分为以下几个主要类别：

GPU 之间的点对点（P2P）传输：允许 CUDA 程序在同一系统中的两个 GPU 之间使用高速直接内存传输（DMA）来复制数据。它还允许对同一系统中其他 GPU 的内存进行优化访问。
网络和存储之间的加速通信：这项技术有助于从第三方设备（如 InfiniBand 网络适配器或存储）直接访问 CUDA 内存。它消除了不必要的内存复制和 CPU 开销，从而减少了传输和访问的延迟。此功能从 CUDA 3.1 开始支持。
视频的 GPUDirect：这项技术优化了基于帧的视频设备的流水线。它允许与 OpenGL、DirectX 或 CUDA 进行低延迟通信，并且从 CUDA 4.2 开始支持。
远程直接内存访问（RDMA）：此功能允许集群中的 GPU 之间进行直接通信。此功能从 CUDA 5.0 及更高版本开始支持。

在本节中，我们将把我们的顺序代码转换为使用 GPUDirect 的 P2P 功能，以便在同一系统中的多个 GPU 上运行。

GPUDirect P2P 功能允许以下操作：

GPUDirect 传输：cudaMemcpy()启动了从 GPU 1 的内存到 GPU 2 的内存的 DMA 复制。
直接访问：GPU 1 可以读取或写入 GPU 2 的内存（加载/存储）。

以下图表展示了这些功能：

要理解 P2P 的优势，有必要了解 PCIe 总线规范。这是为了通过 InfiniBand 等互连优化与其他节点进行通信而创建的。当我们想要从单个 GPU 优化地发送和接收数据时，情况就不同了。以下是一个样本 PCIe 拓扑，其中八个 GPU 连接到各种 CPU 和 NIC/InfiniBand 卡：

在前面的图表中，GPU0 和 GPU1 之间允许 P2P 传输，因为它们都位于同一个 PCIe 交换机中。然而，GPU0 和 GPU4 不能执行 P2P 传输，因为两个I/O Hub（IOHs）之间不支持 PCIe P2P 通信。IOH 不支持来自 PCI Express 的非连续字节进行远程对等 MMIO 事务。连接两个 CPU 的 QPI 链路的性质确保了如果 GPU 位于不同的 PCIe 域上，则不可能在 GPU 内存之间进行直接 P2P 复制。因此，从 GPU0 的内存到 GPU4 的内存的复制需要通过 PCIe 链路复制到连接到 CPU0 的内存，然后通过 QPI 链路传输到 CPU1，并再次通过 PCIe 传输到 GPU4。正如你所想象的那样，这个过程增加了大量的开销，无论是延迟还是带宽方面。

以下图表显示了另一个系统，其中 GPU 通过支持 P2P 传输的 NVLink 互连相互连接：

前面的图表显示了一个样本 NVLink 拓扑，形成了一个八立方网格，其中每个 GPU 与另一个 GPU 最多相连 1 跳。

更重要的问题是，我们如何找出这个拓扑结构以及哪些 GPU 支持 P2P 传输？幸运的是，有工具可以做到这一点。nvidia-smi就是其中之一，它作为 NVIDIA 驱动程序安装的一部分被安装。以下屏幕截图显示了在前面图表中显示的 NVIDIA DGX 服务器上运行nvidia-smi的输出：

前面的屏幕截图代表了在具有 8 个 GPU 的 DGX 系统上运行nvidia-smi topo -m命令的结果。如您所见，通过 SMP 互连（QPI/UPI）连接到另一个 GPU 的任何 GPU 都无法执行 P2P 传输。例如，GPU0将无法与GPU5、GPU6和GPU7进行 P2P 传输。另一种方法是通过 CUDA API 来找出这种传输，我们将在下一节中使用它来转换我们的代码。

现在我们已经了解了系统拓扑，我们可以开始将我们的应用程序转换为单个节点/服务器上的多个 GPU。

单节点-多 GPU 高斯消元

准备您的多 GPU 应用程序。此代码可以在本书的 GitHub 存储库中的06_multigpu/gaussian中找到。使用nvcc编译器编译您的应用程序，如下所示：

$ nvcc -o gaussian_multi_gpu_p2p.out gaussian_multi_gpu_p2p.cu
$ time ./gaussian_multi_gpu_p2p.out

从单 GPU 实现转换为多 GPU 实现，我们在上一小节中定义的三个内核将被原样使用。但是，线性系统被分成与 GPU 数量相等的部分。这些部分分配给每个 GPU 一个部分。每个 GPU 负责对分配给该 GPU 的部分执行操作。矩阵是按列分割的。这意味着每个 GPU 从所有行中获得相等数量的连续列。用于找到主元的内核在包含主元素的列上启动。主元元素的行索引被广播到其他 GPU。提取的主元行和行消除内核在所有 GPU 上启动，每个 GPU 都在矩阵的自己的部分上工作。以下图显示了行在多个 GPU 之间的分割以及主元行需要广播到其他进程的情况：

上述图表示了在多个 GPU 上的工作分配。目前，主元行属于GPU1，负责将主元行广播到其他 GPU。

让我们试着理解这些代码更改，以及用于启用 P2P 功能的 CUDA API：

在支持的 GPU 之间启用 P2P 访问。以下代码显示了这个步骤的第一步：启用 GPU 之间的 P2P 访问：

for( int i = 0; i < nGPUs; i++ ){   
    // setup P2P 
    cudaSetDevice(i);   
    for( int j = 0; j < nGPUs; j++ ) {      
        if (i == j) continue;      
        cudaDeviceCanAccessPeer(&canAccessPeer, i, j);
        if (canAccessPeer)      
            cudaDeviceEnablePeerAccess(j, 0);    
    } 
}

在上述代码中使用的关键 API 如下：

- cudaDeviceCanAccessPeer(): 检查当前 GPU 是否可以对传递的 GPU ID 进行 P2P 访问
cudaDeviceEnablePeerAccess(): 如果cudaDeviceCanAccessPeer()返回True，则启用 P2P 访问

拆分并将内容传输到各自的 GPU：

for( int g = 0; g < nGPUs; g++ ){       
    cudaSetDevice(g);       
    //Copy  part ‘g’ of ABT to GPU ‘g’; 
}

在上述代码中使用的关键 API 是cudaSetDevice()。这将当前上下文设置为作为参数传递的 GPU ID。

找到主元行并通过 P2P 进行广播：

for( int n = 0; n < N; n++ ){        
    gp = GPU that holds n;        
    cudaSetDevice(gp);        
    findPivotRowAndMultipliers<<<...>>>();
    for( int g = 0; g < nGPUs; g++ ){ 
        if (g == gp) continue;
        cudaMemcpyPeer(pivotDatag, g, pivotDatagp, gp, numBytes);
     }  ...

用于将传输广播到 GPU 的 API 是cudaMemcpyPeer()。

提取主元行并执行行消除：

for( int n = 0; n < N; n++ ){
    ...
    for( int g = 0; g < nGPUs; g++ ){  
        cudaSetDevice(g); 
        extractPivotRow<<<...>>>(); 
        rowElimination<<<...>>>();   
    }  
}

如您所见，我们仍在重用相同的内核。唯一的区别是我们使用cudaSetDevice() API 告诉 CUDA 运行时内核应该在哪个 GPU 上启动。请注意，cudaSetDevice()是一个昂贵的调用，特别是在旧一代的 GPU 上。因此，建议您通过在 CPU 上并行调用nGPUs的 for 循环，利用OpenMP/OpenACC或 CPU 上的任何其他线程机制来调用。

从各自的 CPU 中复制数据回来：

for( int g = 0; g < nGPUs; g++ ){ 
    cudaSetDevice(g);  
    Copy  part ‘g’ of reduced ABT from GPU ‘g’ to Host; 
}

这五个步骤完成了将单个 GPU 实现转换为单个节点上的多个 GPU 的练习。

作为 CUDA 安装的一部分提供的 CUDA 示例包括一些测试 P2P 带宽性能的示例代码。它可以在samples/1_Utilities/p2pBandwidthLatencyTest文件夹中找到。建议您在系统上运行此应用程序，以便了解系统的 P2P 带宽和延迟。

现在我们已经在单个节点上实现了多 GPU，我们将改变方向并在多个 GPU 上运行此代码。但在将我们的代码转换为多个 GPU 之前，我们将提供一个关于 MPI 编程的简短介绍，这主要用于节点间通信。

MPI 的简要介绍

消息传递接口（MPI）标准是一种消息传递库标准，已成为在 HPC 平台上编写消息传递程序的行业标准。基本上，MPI 用于在多个 MPI 进程之间进行消息传递。相互通信的 MPI 进程可以驻留在同一节点上，也可以跨多个节点。

以下是一个 Hello World MPI 程序的示例：

#include <mpi.h> 
int main(int argc, char *argv[]) {     
    int rank,size;     
    /* Initialize the MPI library */     
    MPI_Init(&argc,&argv);     
    /* Determine the calling process rank and total number of ranks */
    MPI_Comm_rank(MPI_COMM_WORLD,&rank);     
    MPI_Comm_size(MPI_COMM_WORLD,&size);     
    /* Compute based on process rank */     
    /* Call MPI routines like MPI_Send, MPI_Recv, ... */     
    ...     
    /* Shutdown MPI library */     
    MPI_Finalize();     
    return 0; 
}

正如您所看到的，MPI 程序涉及的一般步骤如下：

我们包括头文件mpi.h，其中包括所有 MPI API 调用的声明。
我们通过调用MPI_Init并将可执行参数传递给它来初始化 MPI 环境。在这个语句之后，多个 MPI 等级被创建并开始并行执行。
所有 MPI 进程并行工作，并使用诸如MPI_Send()、MPI_Recv()等消息传递 API 进行通信。
最后，我们通过调用MPI_Finalize()终止 MPI 环境。

我们可以使用不同的 MPI 实现库（如 OpenMPI、MVPICH、Intel MPI 等）来编译此代码：

$ mpicc -o helloWorldMPI helloWorldMPI.c
$ mpirun -n 4 --hostfile hostsList ./helloWorldMPI

我们使用mpicc编译器来编译我们的代码。mpicc基本上是一个包装脚本，它在内部扩展编译指令，以包括相关库和头文件的路径。此外，运行 MPI 可执行文件需要将其作为参数传递给mpirun。mpirun是一个包装器，它帮助在应用程序应该执行的多个节点上设置环境。-n 4参数表示我们要运行四个进程，并且这些进程将在主机名存储在文件主机列表中的节点上运行。

在本章中，我们的目标是将 GPU 内核与 MPI 集成，使其在多个 MPI 进程中运行。但我们不会涵盖 MPI 编程的细节。那些不熟悉 MPI 编程的人应该先查看computing.llnl.gov/tutorials/mpi/，了解分布式并行编程，然后再进入下一节。

GPUDirect RDMA

在集群环境中，我们希望在多个节点上利用 GPU。我们将允许我们的并行求解器将 CUDA 代码与 MPI 集成，以利用多节点、多 GPU 系统上的多级并行性。使用 CUDA-aware MPI 来利用 GPUDirect RDMA 进行优化的节点间通信。

GPUDirect RDMA 允许在集群中的 GPU 之间进行直接通信。它首先由 CUDA 5.0 与 Kepler GPU 卡支持。在下图中，我们可以看到 GPUDirect RDMA，即Server 1中的GPU 2直接与Server 2中的GPU 1通信：

GPUDirect RDMA 工作的唯一理论要求是网络卡和GPU共享相同的根复杂性。 GPU 和网络适配器之间的路径决定了是否支持 RDMA。让我们重新访问我们在上一节中运行的 DGX 系统上nvidia-smi topo -m命令的输出：

如果我们看一下GPU4行，它显示GPU4到mlx5_2连接类型为PIX（通过 PCIe 交换机遍历）。我们还可以看到GPU4到mlx_5_0连接类型为SYS（通过QPI遍历）。这意味着GPU4可以通过 Mellanox InfiniBand 适配器mlx_5_2执行 RDMA 传输，但如果需要从mlx_5_0进行传输，则无法进行 RDMA 协议，因为QPI不允许。

CUDA-aware MPI

所有最新版本的 MPI 库都支持 GPUDirect 功能。支持 NVIDIA GPUDirect 和统一虚拟寻址（UVA）的 MPI 库使以下功能可用：

MPI 可以将 API 传输直接复制到/从 GPU 内存（RDMA）。
MPI 库还可以区分设备内存和主机内存，无需用户提示，因此对 MPI 程序员透明。
程序员的生产率提高了，因为少量应用代码需要更改以在多个 MPI 秩之间传输数据。

正如我们之前提到的，CPU 内存和 GPU 内存是不同的。没有 CUDA-aware MPI，开发人员只能将指向 CPU/主机内存的指针传递给 MPI 调用。以下代码是使用非 CUDA-aware MPI 调用的示例：

 //MPI rank 0:Passing s_buf residing in GPU memory 
 // requires it to be transferred to CPU memory
cudaMemcpy(s_buf_h,s_buf_d,size,cudaMemcpyDeviceToHost);
MPI_Send(s_buf_h,size,MPI_CHAR,1,100,MPI_COMM_WORLD);

//MPI rank 1: r_buf received buffer needs to be 
// transferred to GPU memory before being used in GPU
MPI_Recv(r_buf_h,size,MPI_CHAR,0,100,MPI_COMM_WORLD, &status);
cudaMemcpy(r_buf_d,r_buf_h,size,cudaMemcpyHostToDevice);

有了 CUDA-aware MPI 库，这是不必要的；GPU 缓冲区可以直接传递给 MPI，如下所示：

//MPI rank 0
MPI_Send(s_buf_d,size,MPI_CHAR,1,100,MPI_COMM_WORLD);

//MPI rank n-1
MPI_Recv(r_buf_d,size,MPI_CHAR,0,100,MPI_COMM_WORLD, &status);

例如，对于 Open MPI，CUDA-aware 支持存在于 Open MPI 1.7 系列及更高版本中。要启用此功能，需要在编译时配置 Open MPI 库以支持 CUDA，如下所示：

$ ./configure --with-cuda

拥有 CUDA-aware MPI 并不意味着总是使用 GPUDirect RDMA。如果数据传输发生在网络卡和 GPU 之间共享相同的根复杂，则使用 GPUDirect 功能。尽管如此，即使未启用 RDMA 支持，拥有 CUDA-aware MPI 也可以通过利用诸如消息传输之类的功能使应用程序更有效，如下图所示可以进行流水线处理：

！

上图显示了具有 GPUDirect 的 CUDA-aware MPI 与不具有 GPUDirect 的 CUDA-aware MPI。两个调用都来自 CUDA-aware MPI，但左侧是 GPUDirect 传输，右侧是没有 GPUDirect 传输。

非 GPUDirect 传输有以下阶段：

节点 1：从 GPU1 传输到主机内存
节点 1：从主机内存传输到网络适配器暂存区
网络：通过网络传输
节点 2：从网络暂存区传输到主机内存
节点 2：从主机内存传输到 GPU 内存

如果支持 GPUDirect RDMA，则从 GPU 传输直接通过网络进行，涉及主机内存的额外副本都被删除。

现在我们已经掌握了这个概念，让我们开始将代码转换为使用 CUDA-aware MPI 编程启用多 GPU 支持。

多节点-多 GPU 高斯消元

准备您的 GPU 应用程序。此代码可以在本书的 GitHub 存储库中的06_multigpu/gaussian中找到。使用nvcc编译器编译和运行应用程序，如下所示：

$ mpicc-o gaussian_multi_gpu_rdma.out gaussian_multi_gpu_rdma.cu
$ mpirun -np 8 ./gaussian_multi_gpu_rdma.out

我们使用mpicc而不是nvcc来编译 MPI 程序。我们使用mpirun命令运行可执行文件，而不是直接运行已编译的可执行文件。本节中您将看到的结果是在同一系统上具有 8 个 V100 的 DGX 系统上运行的输出。我们利用 8 个最大 MPI 进程，将每个 GPU 映射为 1 个 MPI 进程。要了解如何将多个 MPI 进程映射到同一 GPU，请阅读本章后面的MPS子节。在本练习中，我们使用了已编译为支持 CUDA 的 Open MPI 1.10，如前一节所述。

多 GPU 实现涉及的步骤如下：

MPI 进程的秩 0 生成线性系统（矩阵 A，B）的数据。
转置增广矩阵（AB^T）由根节点在 MPI 进程之间使用MPI_Scatterv()按行分割。
每个 MPI 进程并行计算其部分输入：

三个内核的处理发生在 GPU 上。
在findPivot操作后，通过MPI_Send()/Recv()实现了枢轴的共识。

减少的转置增广矩阵（ABT）使用MPI_Gatherv()在根节点上收集。
根节点执行回代以计算解 X。

展示前面代码的提取样本高斯代码如下：

void gaussianEliminationOnGPU() {
    cudaSetDevice(nodeLocalRank); //Set CUDA Device based on local rank
    //Copy  chuck of AB Transpose from Host to GPU; 
   for( int n = 0; n < N; n++ ){ 
       prank = MPI rank that holds n; 
       if (myRank == prank) 
           findPivotRowAndMultipliers<<<...>>>(); 
       bCastPivotInfo(); // from prank to other ranks 
       extractPivotRow<<<...>>>(); 
       rowElimination<<<...>>>(); 
   //Copy  myPartOfReducedTransposeAB from GPU to Host;
}

现在，让我们添加多 GPU 支持：

设置每个 MPI 等级的 CUDA 设备：在 Open MPI 中，您可以通过使用MPI_COMM_TYPE_SHARED作为MPI_Comm_split_type的参数来获得 MPI 进程的本地等级，如下面的代码所示：

MPI_Comm loc_comm;
MPI_Comm_split_type(MPI_COMM_WORLD, MPI_COMM_TYPE_SHARED, rank, MPI_INFO_NULL, &loc_comm);
int local_rank = -1;
MPI_Comm_rank(loc_comm,&local_rank);
MPI_Comm_free(&loc_comm);

现在我们有了本地等级，每个 MPI 进程都使用它来通过cudaSetDevice()设置当前 GPU，如下图所示：

使用MPI_Scatter将输入拆分并分发到不同的 MPI 进程：

void distributeInputs() {
    MPI_Scatterv(transposeAB, ..., myPartOfTransposeAB, recvCount, MPI_UNSIGNED, 0, MPI_COMM_WORLD); 
}

在 GPU 上执行高斯消元：

void gaussianEliminationOnGPU() { 
    cudaSetDevice(nodeLocalRank);
     for( int n = 0; n < N; n++ ){ 
        prank = MPI rank that holds n; 
        if (myRank == prank) 
            findPivotRowAndMultipliers<<<...>>>();
        MPI_Bcast(...); // from prank to other ranks 
        extractPivotRow<<<...>>>(); 
        rowElimination<<<...>>>(); 
}

在执行任何操作之前，基于本地等级设置当前 GPU。然后，由负责该行的进程提取枢轴行，然后将枢轴行广播到所有其他 MPI 等级，我们用于消除。

通过使用异步 MPI 调用而不是使用广播 API（如MPI_Bcast），可以提高传输时间的整体性能。实际上，不建议使用广播 API；它应该被替换为可以实现相同功能的MPI_Isend和MPI_Irecv，这些是异步版本。请注意，使调用异步会增加其他方面（如调试）的复杂性。因此，用户需要编写额外的代码来发送和接收数据。

本章提供了在向现有 MPI 程序添加 GPU 支持时的最佳编码实践，并不应被视为 MPI 编程的最佳编程实践的专家指南。

CUDA 流

流以 FIFO 方式工作，其中操作的顺序按照它们发出的顺序执行。从主机代码发出的请求被放入先进先出队列中。队列由驱动程序异步读取和处理，并且设备驱动程序确保队列中的命令按顺序处理。例如，内存复制在内核启动之前结束，依此类推。

使用多个流的一般想法是，在不同流中触发的 CUDA 操作可能会并发运行。这可能导致多个内核重叠或内核执行中的内存复制重叠。

为了理解 CUDA 流，我们将看两个应用程序。第一个应用程序是一个简单的矢量加法代码，添加了流，以便它可以重叠数据传输和内核执行。第二个应用程序是一个图像合并应用程序，也将在第九章中使用，使用 OpenACC 进行 GPU 编程。

首先，根据以下步骤配置您的环境：

准备您的 GPU 应用程序。例如，我们将合并两个图像。此代码可以在本书的 GitHub 存储库的06_multi-gpu/streams文件夹中找到。
使用nvcc编译器编译您的应用程序如下：

$ nvcc --default-stream per-thread -o vector_addition -Xcompiler -fopenmp -lgomp vector_addition.cu
$ nvcc --default-stream per-thread -o merging_muli_gpu -Xcompiler -fopenmp -lgomp scrImagePgmPpmPackage.cu image_merging.cu
$ ./vector addition
$ ./merging_muli_gpu

上述命令将创建两个名为vector_addition和merging_multi_gpu的二进制文件。正如您可能已经注意到的，我们在我们的代码中使用了额外的参数。让我们更详细地了解它们：

--default-stream per-thread：此标志告诉编译器解析代码中提供的 OpenACC 指令。
-Xcompiler -fopenmp -lgomp：此标志告诉nvcc将这些附加标志传递给 CPU 编译器，以编译代码的 CPU 部分。在这种情况下，我们要求编译器向我们的应用程序添加与 OpenMP 相关的库。

我们将把这一部分分为两部分。应用程序 1 和应用程序 2 分别演示了在单个和多个 GPU 中使用流。

应用程序 1-使用多个流来重叠数据传输和内核执行

我们需要遵循的步骤来重叠数据传输和内核执行，或者同时启动多个内核如下：

声明要固定的主机内存，如下面的代码片段所示：

cudaMallocHost(&hostInput1, inputLength*sizeof(float));
cudaMallocHost(&hostInput2, inputLength*sizeof(float));
cudaMallocHost(&hostOutput, inputLength*sizeof(float));

在这里，我们使用cudaMallocHost() API 来分配固定内存的向量。

创建一个Stream对象，如下面的代码片段所示：

for (i = 0; i < 4; i++) {
 cudaStreamCreateWithFlags(&stream[i],cudaStreamNonBlocking);

在这里，我们使用cudaStreamCreateWithFlags() API，传递cudaStreamNonBlocking作为标志，使此流非阻塞。

调用 CUDA 内核和内存复制时使用stream标志，如下面的代码片段所示：

for (i = 0; i < inputLength; i += Seglen * 4) {
    for (k = 0; k < 4; k++) {
        cudaMemcpyAsync(... , cudaMemcpyHostToDevice, stream[k]);
        cudaMemcpyAsync(... , cudaMemcpyHostToDevice, stream[k]);
        vecAdd<<<Gridlen, 256, 0, stream[k]>>>(...);
    }
}

如我们所见，我们不是通过一次复制整个数组来执行矢量加法，而是将数组分成段，并异步复制这些段。内核执行也是在各自的流中异步进行的。

当我们通过 Visual Profiler 运行这段代码时，我们可以看到以下特点：

前面的分析器截图显示，蓝色条（基本上是vector_addition内核）重叠了内存复制。由于我们在代码中创建了四个流，分析器中也有四个流。

每个 GPU 都有两个内存复制引擎。一个负责主机到设备的传输，另一个负责设备到主机的传输。因此，发生在相反方向的两个内存复制可以重叠。此外，内存复制可以与计算内核重叠。这可以导致n路并发，如下图所示：

每个 GPU 架构都有一定的约束和规则，根据这些规则，我们将在执行时看到这些重叠。一般来说，以下是一些指导方针：

CUDA 操作必须在不同的非 0 流中。
使用cudaMemcpyAsync时，主机应该使用cudaMallocHost()或cudaHostAlloc()进行固定。
必须有足够的资源可用。
不同方向的cudaMemcpyAsyncs
设备资源（SMEM、寄存器、块等）以启动多个并发内核

应用程序 2 - 使用多个流在多个设备上运行内核

为了在多个设备上运行内核并重叠内存传输，我们之前遵循的步骤保持不变，除了一个额外的步骤：设置 CUDA 设备以创建流。让我们看看以下步骤：

创建与系统中 CUDA 设备数量相等的流，如下面的代码片段所示：

cudaGetDeviceCount(&noDevices);
cudaStream_t *streams;
streams = (cudaStream_t*) malloc(sizeof(cudaStream_t) * noDevices);

我们使用cudaGetDeviceCount() API 来获取 CUDA 设备的数量。

在各自的设备中创建流，如下面的代码片段所示：

#pragma omp parallel num_threads(noDevices)
{
     int block = omp_get_thread_num();
    cudaSetDevice(block);
    cudaStreamCreate(&streams[block]);

我们启动与 CUDA 设备数量相等的 OpenMP 线程，以便每个 CPU 线程可以为其各自的设备创建自己的 CUDA 流。每个 CPU 线程执行cudaSetDevice()来根据其 ID 设置当前 GPU，然后为该设备创建流。

在该流中启动内核和内存复制，如下所示：

cudaMemcpyAsync(... cudaMemcpyHostToDevice,streams[block]);
cudaMemcpyAsync(..., cudaMemcpyHostToDevice, streams[block]);
merging_kernel<<<gridDim,blockDim,0,streams[block]>>>(...);
cudaMemcpyAsync(...,streams[block]);

在分析器中运行代码后的输出可以在下面的截图中看到，这代表了 Visual Profiler 的时间轴视图。这显示了一个 GPU 的内存复制与另一个 GPU 的内核执行重叠：

如您所见，我们在拥有四个 V100 的多 GPU 系统上运行了这段代码。不同 GPU 中的内存复制和内核重叠。在这段代码中，我们演示了利用 OpenMP 在不同设备上并行调用 CUDA 内核。这也可以通过利用 MPI 来启动利用不同 GPU 的多个进程来实现。

在下一节中，我们将看一些额外的主题，这些主题可以提高多 GPU 应用程序的性能，并帮助开发人员分析和调试他们的代码。

额外的技巧

在本节中，我们将涵盖一些额外的主题，这些主题将帮助我们了解多 GPU 系统的额外特性。

使用 InfiniBand 网络卡对现有系统进行基准测试

有不同的基准可用于测试 RDMA 功能。InfiniBand 适配器的一个这样的基准可以在www.openfabrics.org/找到。您可以通过执行以下代码来测试您的带宽：

$ git clone git://git.openfabrics.org/~grockah/perftest.git
$ cd perftest 
$ ./autogen.sh 
$ export CUDA_H_PATH=<<Path to cuda.h>> 
$ ./configure –prefix=$HOME/test 
$ make all install

然后，您可以运行以下命令来测试带宽：

For example host to GPU memory (H-G) BW test:
server$ ~/test/bin/ib_write_bw -n 1000 -O -a --use_cuda
client $ ~/test/bin/ib_write_bw -n 1000 -O -a server.name.org

//GPU to GPU memory (G-G) BW test:
server$ ~/test/bin/ib_write_bw -n 1000 -O -a --use_cuda
client $ ~/test/bin/ib_write_bw -n 1000 -O -a --use_cuda server.name.org

NVIDIA 集体通信库（NCCL）

NCCL 提供了常用于深度学习等领域的通信原语的实现。NCCL 1.0 从同一节点内多个 GPU 之间的通信原语实现开始，并发展到支持多个节点上的多个 GPU。NCCL 库的一些关键特性包括以下内容：

支持来自多个线程和多个进程的调用
支持多个环和树拓扑，以更好地利用节点内和节点间的总线
支持 InfiniBand 节点间通信
源代码包可以从 GitHub（github.com/nvidia/nccl）免费下载

NCCL 可以扩展到 24,000 个 GPU，延迟低于 300 微秒。请注意，尽管 NCCL 已被证明是深度学习框架中非常有用和方便的库，但在用于 HPC 应用时存在局限，因为它不支持点对点通信。NCCL 支持集体操作，这在深度学习应用中被使用，例如以下内容：

AllReduce
AllGather
ReduceScatter
Reduce
Broadcast

所有 NCCL 调用都作为 CUDA 内核运行，以更快地访问 GPU 内存。它使用较少的线程，实现为一个块。这最终只在一个 GPU SM 上运行，因此不会影响其他 GPU 的利用率。让我们看一下以下代码：

ncclGroupStart(); 
for (int i=0; i<ngpus; i++) 
{ 
    ncclAllGather(…, comms[i], streams[i]); 
} 
ncclGroupEnd();

正如我们所看到的，NCCL 调用简单，易于调用。

使用 NCCL 加速集体通信

NVIDIA 集体通信库（NCCL）提供了为多个 NVIDIA GPU 优化的性能集体通信原语。在本节中，我们将看到这个库是如何工作的，以及我们如何从中受益。

并不难找到使用多个 GPU 来训练网络的深度学习模型。由于两个 GPU 并行计算神经网络，我们很容易想象这种技术将随着 GPU 数量的增加而提高训练性能。不幸的是，世界并不那么简单。梯度应该在多个 GPU 之间共享，并且一个 GPU 中的权重更新过程应该等待其他 GPU 的梯度来更新其权重。这是使用多个 GPU 进行深度学习训练的一般过程，并在以下图表中显示：

集体通信有许多类型：全局归约、广播、归约、全局收集、归约散射等。在深度学习中，每个 GPU 在传输自己的数据的同时收集另一个 GPU 的数据。因此，我们可以确定深度学习在通信中需要所有类型的归约样式通信。

在 HPC 社区中，包括全局归约在内的集体通信是一个常见的话题。节点内和节点间处理器之间的通信是一个具有挑战性但至关重要的问题，因为它直接关系到可扩展性。正如我们在第六章中提到的，可扩展的多 GPU 编程，在多 GPU 编程部分，需要仔细考虑与每个 GPU 的通信。开发人员应该设计和实现 GPU 中的集体通信，即使 MPI 已经支持这样的通信模式。

NCCL 提供了一种集体通信，它了解 GPU 拓扑配置。通过使用各种分组和通信命令，您可以应用所需的通信任务。

一个前提是您的系统需要有多个 GPU，因为 NCCL 是一个与多个 GPU 一起工作的通信库。

以下步骤涵盖了如何调用ncclAllReduce()来测试和测量系统的 GPU 网络带宽。示例代码实现在04_nccl中：

让我们定义一个类型，它将包含、发送和接收每个 GPU 设备的缓冲区和cudaStream，如下所示：

typedef struct device
{
    float *d_send;
    float *d_recv;
    cudaStream_t stream;
} device_t;

在应用程序开始时，我们需要准备一些句柄，以便我们可以控制多个 GPU：

cudaGetDeviceCount(&num_dev);
ncclComm_t *ls_comms = new ncclComm_t[num_dev];
int *dev_ids = new int[num_dev];
for (int i = 0; i < num_dev; i++)
    dev_ids[i] = i;

然后，我们将创建一个缓冲区，假设我们有数据。对于每个设备，我们将初始化每个设备的项目，如下所示：

unsigned long long size = 512 * 1024 * 1024; // 2 GB

// allocate device buffers and initialize device handles
device_t *ls_dev = new device_t[num_dev];
for (int i = 0; i < num_dev; i++) {
    cudaSetDevice(i);
    cudaMalloc((void**)&ls_dev[i].d_send, sizeof(float) * size);
    cudaMalloc((void**)&ls_dev[i].d_recv, sizeof(float) * size);
    cudaMemset(ls_dev[i].d_send, 0, sizeof(float) * size);
    cudaMemset(ls_dev[i].d_recv, 0, sizeof(float) * size);
    cudaStreamCreate(&ls_dev[i].stream);
}

在开始 NCCL 通信之前，我们需要初始化 GPU 设备，以便它们知道它们在 GPU 组中的排名。由于我们将用单个进程测试带宽，我们可以安全地调用一个初始化所有设备的函数：

ncclCommInitAll(ls_comms, num_dev, dev_ids);

如果我们要用多个进程测试带宽，我们需要调用ncclCommInitRank()。我们需要为计算进程 ID 和 GPU 排名提供 GPU ID。
现在，我们可以使用 NCCL 完成 all-reduce 操作。以下代码是ncclAllReduce的示例实现：

ncclGroupStart();
for (int i = 0; i < num_dev; i++) {
    ncclAllReduce((const void*)ls_dev[i].d_send, 
                  (void*)ls_dev[i].d_recv,
        test_size, ncclFloat, ncclSum, 
        ls_comms[i], ls_dev[i].stream);
}
ncclGroupEnd();

对于每个设备，我们需要触发流量。为此，我们需要启动和关闭 NCCL 组通信。现在，我们已经实现了一些使用ncclAllReduce()的测试代码。让我们通过微基准测试来了解 NCCL 的工作原理。

在多 GPU 系统上测试此代码，运行以下命令：

$ nvcc -run -m64 -std=c++11 -I/usr/local/cuda/samples/common/inc -gencode arch=compute_70,code=sm_70 -lnccl -o nccl ./nccl.cu

以下图表显示了在 DGX Station 中使用四个 V100 32G GPU 测得的性能。蓝线表示基于 NVLink 的带宽，而橙线表示基于 PCIe 的带宽，通过设置NCCL_P2P_DISABLE=1 ./ncd并关闭对等 GPU 来实现：

这个 NCCL 测试可能会受到系统配置的影响。这意味着结果可能会有所不同，取决于您系统的 GPU 拓扑结构。

这显示了基于 PCI Express 和基于 NVLINK 的 all-reduce 性能差异。我们可以使用nvprof来查看通信。以下屏幕截图显示了通过 NCCL 2.3.7 在 DGX Station 上的 all-reduce 通信：

NCCL 越来越快。通过引入新的 GPU 互连技术 NVLink 和 NVSwitch，我们对 NCCL 的经验正在增加，以至于我们可以实现可扩展的性能。

以下链接提供了关于 NCCL 的讨论：developer.nvidia.com/gtc/2019/video/S9656/video。

摘要

在本章中，我们介绍了多 GPU 编程的不同方法。通过示例高斯消元，我们看到了如何将单个 GPU 应用程序工作负载分割到多个 GPU 中，首先是单个节点，然后是多个节点。我们看到了系统拓扑在利用 P2P 传输和 GPUDirect RDMA 等功能方面起着重要作用。我们还看到了如何使用多个 CUDA 流来重叠多个 GPU 之间的通信和数据传输。我们还简要介绍了一些其他主题，可以帮助 CUDA 程序员优化代码，如 MPS 和使用nvprof来分析多 GPU 应用程序。

在下一章中，我们将看到大多数 HPC 应用程序中出现的常见模式以及如何在 GPU 中实现它们。

第七章：CUDA 中的并行编程模式

在本章中，我们将涵盖并行编程算法，这将帮助您了解如何并行化不同的算法并优化 CUDA。本章中我们将涵盖的技术可以应用于各种问题，例如我们在第三章中看到的并行减少问题，CUDA 线程编程，它可以用于设计神经网络操作中的高效 softmax 层。

在本章中，我们将涵盖以下主题：

矩阵乘法优化
图像卷积
前缀和
打包和拆分
N 体操作
在 CUDA 中使用动态并行性进行快速排序
基数排序
直方图计算

技术要求

为了完成本章，建议您使用 Pascal 架构之后的 NVIDIA GPU 卡。换句话说，您的 GPU 的计算能力应等于或大于 60。如果您不确定您的 GPU 的架构，请访问 NVIDIA GPU 的网站（developer.nvidia.com/cuda-gpus）并确认您的 GPU 的计算能力。

本章中的相同代码已经使用 CUDA 版本 10.1 进行开发和测试。一般来说，如果适用的话，建议使用最新的 CUDA 版本。

矩阵乘法优化

虽然我们在许多示例中使用了矩阵乘法代码，但我们并没有调查操作是否被优化。现在，让我们回顾其操作以及如何找到优化的机会。

矩阵乘法是从两个矩阵进行的一组点积运算。我们可以简单地并行化所有 CUDA 线程执行的操作，以生成元素的点积。然而，从内存使用的角度来看，这种操作效率低，因为从内存加载的数据没有被重复使用。为了确认我们的类比，让我们测量性能限制器。以下图表显示了使用 NVIDIA Nsight Compute 的 Tesla V100 卡的 GPU 利用率：

根据我们的性能限制器分析，这种利用率可以归类为内存受限。因此，我们应该审查内存利用率以减少利用率。以下截图显示了内存工作负载分析部分：

通过这个分析，我们可以看到 L2 缓存命中率低，最大带宽也低。我们可以推测这是因为原始矩阵乘法操作没有重复使用加载的数据，正如我们之前提到的。这可以通过使用共享内存来解决，即重复使用加载的数据并减少全局内存使用。现在，让我们回顾矩阵乘法以及如何优化使用具有小内存空间的共享内存。

矩阵乘法是一组点积运算，使用一些小尺寸矩阵和输出的累积。小矩阵称为瓦片，它们映射到输出矩阵上。每个瓦片将并行计算自己的输出。这个操作可以按以下步骤实现：

确定两个输入和输出矩阵的瓦片大小。
遍历输入瓦片，以及它们的方向（矩阵 A 向右移动，矩阵 B 向下移动）。
在瓦片内计算矩阵乘法。
继续第二步，直到瓦片达到末尾。
刷新输出。

以下图表显示了瓦片矩阵乘法的概念：

在上图中，我们计算矩阵乘法，C = AB。我们从矩阵 A 和矩阵 B 中计算一个较小的矩阵乘法作为瓦片（绿色）。然后，我们分别遍历输入瓦片位置。操作结果累积到先前的输出，以生成矩阵乘法的输出。

这个操作提供了一个优化机会，因为我们可以将大矩阵操作分解为小问题，并将其放置在小内存空间中。在 CUDA 编程中，我们将小矩阵放置在共享内存中，并减少全局内存访问。在我们的实现中，我们将瓦片与 CUDA 线程块匹配。瓦片的位置将由其块索引确定，这是通过tid_*变量完成的。

实现平铺方法

现在，让我们使用平铺方法实现优化的矩阵乘法。我们将重用之前在第三章中使用的矩阵乘法示例代码，即 CUDA 线程编程。优化后，我们将看看如何提高性能。按照以下步骤开始：

让我们创建一个核函数，这将是我们优化版本的矩阵乘法。我们将在sgemm操作中命名核函数为v2。这个核函数将计算，因此我们应该分别提供相关参数。我们还将使用M、N和K传递矩阵大小信息：

__global__ void sgemm_kernel_v2(const float *A, const float *B, float *C,
    int M, int N, int K, float alpha, float beta) {}

对于这个操作，我们将分别使用块索引和线程索引。正如我们之前讨论的，我们需要单独使用块索引来指定瓦片位置。我们将使用线程索引进行瓦片级矩阵乘法。因此，我们需要创建 CUDA 索引参数，如下所示：

int bid_x = blockIdx.x * blockDim.x;
int bid_y = blockIdx.y * blockDim.y;
int tid_x = threadIdx.x;
int tid_y = threadIdx.y;

之后，我们将使用共享内存作为瓦片，并使用本地寄存器保存输出值：

float element_c = 0.f;
__shared__ float s_tile_A[BLOCK_DIM][BLOCK_DIM];
__shared__ float s_tile_B[BLOCK_DIM][BLOCK_DIM];

然后，我们将编写一个控制瓦片位置的循环。以下是基于其块大小控制循环的 for 循环代码。请注意，循环大小由K决定，考虑到块应该遍历多少次：

for (int k = 0; k < K; k += BLOCK_DIM)
{
   ... {step 5 and 6 will cover } ...
}

现在，我们将编写代码，将数据输入第二个循环。正如我们之前讨论的，每个瓦片都有自己的移动方向，以及矩阵；瓦片A遍历矩阵A的列，瓦片B遍历矩阵B的行。我们根据矩阵乘法优化部分中显示的图表来放置它们。之后，我们应该在从全局内存复制数据到共享内存后放置__syncthreads()，以避免来自上一次迭代的未更新数据：

// Get sub-matrix from A
s_tile_A[tid_y][tid_x] = A[ (bid_y + tid_y) * K + tid_x + k ];
// Get sub-matrix from B 
s_tile_B[tid_y][tid_x] = B[ k * N + bid_x + tid_x ]; 

__syncthreads();

然后，我们可以从瓦片中编写矩阵乘法代码。名为element_c的本地变量将累积结果：

for (int e = 0; e < BLOCK_DIM; e++)
    element_c += s_tile_A[tid_y][e] * s_tile_B[e][tid_x];

我们将结果写入全局内存。以下操作应该放置在第二个循环完成后：

C[(bid_y + tid_y) * N + (bid_x + tid_x)] = \
 alpha * element_c + beta * C[(bid_y + tid_y) * N + (bid_x + tid_x)];

现在，让我们回顾一下这种平铺方法如何有利于矩阵乘法操作。通过在我们的平铺矩阵乘法中使用共享内存，我们可以期望通过使用输入数据减少全局内存流量，从而增强性能。我们可以轻松地通过配置文件结果来确认这一点：

$ nvcc -m64 -I/usr/local/cuda/samples/common/inc -gencode arch=compute_70,code=sm_70 -o sgemm ./sgemm.cu 
$ nvprof ./sgemm 

        Type Time(%)    Time Calls      Avg      Min      Max Name
GPU activities: 47.79% 9.9691ms     1 9.9691ms 9.9691ms 9.9691ms sgemm_kernel(...)
 32.52% 6.7845ms     1 6.7845ms 6.7845ms 6.7845ms sgemm_kernel_v2(...)

由于我们设计了核心以重用输入数据，增加的块大小可能有助于性能。例如，考虑到 warp 大小和共享内存银行的数量，32 x 32 的块大小可能是最佳的，以避免银行冲突。我们可以轻松地使用配置文件获得其实验结果：

 Type Time(%)    Time Calls      Avg       Min       Max Name
GPU activities: 46.52% 8.1985ms     1 8.1985ms  8.1985ms  8.1985ms sgemm_kernel(...)
 31.24% 5.4787ms     1 5.4787ms  5.4787ms  5.4787ms sgemm_kernel_v2(...)

正如你所看到的，增加的瓦片大小有利于矩阵乘法操作的性能。现在，让我们分析其性能。

平铺方法的性能分析

之前，我们看过了平铺方法以及它如何能够实现良好的性能。让我们回顾一下平铺方法解决了什么问题，并看看接下来我们可以采取哪些步骤。总的来说，覆盖这部分是可选的，因为 NVIDIA 提供了 cuBLAS 和 CUTLASS 库，用于提供优化性能的 GEMM（通用矩阵乘法）操作。

下图显示了来自 NVIDIA Nsight Compute 的更新后的 GPU 利用率报告。较低配置文件的更新利用率输出是上配置文件的结果：

由于两个资源都利用率很高，我们应该审查每个资源的资源使用情况。首先，让我们来审查内存工作量。以下截图显示了更新后的结果：

从这个结果可以看出，全局内存访问从最大化内存带宽和减少内存吞吐量进行了优化。此外，L2 缓存命中率也得到了提高。因此，我们的平铺方法将矩阵乘法从全局内存转换为芯片级操作。

然而，这并不意味着我们已经实现了最优化的性能。从内存工作量分析中，我们可以看到内存管道太忙了。这是由于我们从共享内存进行逐元素乘法。为了解决这个问题，我们需要重新映射共享内存中的数据。我们不会在这本书中涉及到这个问题，但你可以在这篇文章中了解到：github.com/NervanaSystems/maxas/wiki/SGEMM。

正如我们之前讨论的，cuBLAS 库显示出更快的性能。我们将在第八章的cuBLAS部分中介绍其用法。然而，在这个阶段理解平铺方法是有用的，这样我们就可以理解 GPU 如何开始优化。

卷积

卷积操作（或滤波）是许多应用中常见的操作，特别是在图像和信号处理以及深度学习中。虽然这个操作是基于输入和滤波器的顺序数据的乘积，但我们对矩阵乘法有不同的方法。

CUDA 中的卷积操作

卷积操作包括源数据和滤波器。滤波器也被称为核。通过将滤波器应用于输入数据，我们可以获得修改后的结果。下图显示了二维卷积的示意图：

当我们实现卷积操作时，我们需要考虑一些概念，即核和填充。核是一组我们想要应用到源数据的系数。这也被称为滤波器。填充是源数据周围的额外虚拟空间，以便我们可以将核函数应用到边缘。当填充大小为 0 时，我们不允许滤波器移动超出源空间。然而，一般来说，填充大小是滤波器大小的一半。

为了轻松开始，我们可以考虑以下几点来设计核函数：

每个 CUDA 线程生成一个滤波输出。
每个 CUDA 线程将滤波器的系数应用于数据。
滤波器的形状是盒状滤波器。

在满足这些条件的情况下，我们可以有一个简单的卷积操作滤波器，如下所示：

__global__ void
convolution_kernel_v1(float *d_output, float *d_input, float *d_filter, int num_row, int num_col, int filter_size)
{
    int idx_x = blockDim.x * blockIdx.x + threadIdx.x;
    int idx_y = blockDim.y * blockIdx.y + threadIdx.y;

    float result = 0.f;
    // iterates over the every value in the filter
    for (int filter_row = -filter_size / 2; 
         filter_row <= filter_size / 2; ++filter_row)
    {
        for (int filter_col = -filter_size / 2; 
             filter_col <= filter_size / 2; ++filter_col)
        {
            // Find the global position to apply the given filter
            // clamp to boundary of the source
            int image_row = min(max(idx_y + filter_row, 0), 
                                static_cast<int>(num_row - 1));
            int image_col = min(max(idx_x + filter_col, 0), 
                                static_cast<int>(num_col - 1));

            float image_value = static_cast<float>(
                                d_input[image_row * num_col + 
                                image_col]);
            float filter_value = d_filter[(filter_row + 
                                           filter_size / 2) * 
                                           filter_size 
                                           + filter_col + 
                                           filter_size / 2];

            result += image_value * filter_value;
        }
    }

    d_output[idx_y * num_col + idx_x] = result;
}

这个核函数获取输入数据和滤波器进行操作，并没有重用所有数据。考虑到内存效率带来的性能影响，我们需要设计我们的核心代码，以便可以重用加载的数据。现在，让我们编写卷积的优化版本。

优化策略

首先，卷积滤波器是一个只读矩阵，并且被所有 CUDA 线程使用。在这种情况下，我们可以使用 CUDA 的常量内存来利用其缓存操作和广播操作。

在卷积实现设计中，我们使用平铺方法，每个平铺将生成映射位置的滤波输出。我们的平铺设计有额外的空间来考虑卷积滤波器的大小，这为卷积操作提供了所需的数据。这个额外的空间被称为填充。下图显示了一个具有 6 x 6 维度和 3 x 3 大小滤波器的线程块的示例。

然后，我们需要为每个线程块在共享内存上有一个 8 x 8 大小的平铺，如下所示：

当源地址无效内存空间时，或者填充为零（零填充方法）时，填充区域可以是输入数据。通过这样做，我们可以使瓷砖替换输入全局内存而不会对边界元素产生额外影响。为了填充瓷砖，我们使用线程块大小迭代瓷砖，并通过检查输入数据的边界条件来确定应该填充哪个值。我们的实现将输入数据设置为瓷砖大小的倍数，以便边界条件与每个线程块的瓷砖的填充空间匹配。将源数据映射到瓷砖的简要图示如下：

在这个设计中，我们需要做的迭代次数来填充瓷砖是四次。然而，这应该根据滤波器大小进行更改。这样，填充瓷砖的迭代次数由瓷砖大小的上限除以线程块大小确定。其实现很简单，如下面的代码所示：

for (int row = 0; row <= tile_size / BLOCK_DIM; row++) {
    for (int col = 0; col <= tile_size / BLOCK_DIM; col++) {
        ... (filter update operation) ...
    }
}

现在，让我们使用共享内存作为盒式滤波器来实现优化的卷积操作。

使用常量内存优化滤波系数

首先，我们将学习如何优化滤波系数数据的使用。

我们将制作convolution_kernel()的修改版本。让我们复制内核代码，并将其中一个重命名为convolution_kernel_v2()：

首先，我们将创建一个常量内存空间来存储滤波系数。常量内存的大小是有限的，我们不能对内核代码进行修改。然而，我们可以使用这个常量内存，因为我们的卷积滤波器适合这种条件。我们可以这样使用常量内存：

#define MAX_FILTER_LENGTH 128
__constant__ float c_filter[MAX_FILTER_LENGTH * MAX_FILTER_LENGTH];

然后，我们可以使用cudaMemcpyToSymbol()函数将卷积滤波系数放置在常量内存中：

cudaMemcpyToSymbol(c_filter, h_filter, filter_size * filter_size * sizeof(float));

让我们切换滤波操作，这样我们就可以使用常量内存。整个内核实现如下。正如你所看到的，只有一个变量的使用发生了变化：

__global__ void
convolution_kernel_v2(float *d_output, float *d_input, float *d_filter, int num_row, int num_col, int filter_size)
{
    int idx_x = blockDim.x * blockIdx.x + threadIdx.x;
    int idx_y = blockDim.y * blockIdx.y + threadIdx.y;

    float result = 0.f;
    for (int filter_row = -filter_size / 2; 
         filter_row <= filter_size / 2; ++filter_row)
    {
        for (int filter_col = -filter_size / 2; 
             filter_col <= filter_size / 2; ++filter_col)
        {
            int image_row = idx_y + filter_row;
            int image_col = idx_x + filter_col;

            float image_value = (image_row >= 0 
                                 && image_row < num_row 
                                 && image_col >= 0
                                 && image_col < num_col) ?
                                 d_input[image_row * num_col 
                                         + image_col] : 0.f;
            float filter_value = c_filter[(filter_row 
                                          + filter_size / 2) 
                                          * filter_size 
                                          + filter_col 
                                          + filter_size / 2];

            result += image_value * filter_value;
        }
    }

    d_output[idx_y * num_col + idx_x] = result;
}

现在，我们可以通过重复使用nvprof来确认性能提升：

$ nvcc -run -m64 -I/usr/local/cuda/samples/common/inc -gencode arch=compute_70,code=sm_70 -o convolution ./convolution.cu
$ nvprof ./convolution
           Type Time(%) Time Calls Avg Min Max Name
 12.85% 442.21us     1 442.21us 442.21us 442.21us convolution_kernel_v1(...)
 11.97% 412.00us     1 412.00us 412.00us 412.00us convolution_kernel_v2(...)

从这个结果中，我们可以看到减少的内核执行时间。

使用共享内存平铺输入数据

现在，我们将使用共享内存来优化输入数据的使用。为了区分我们的下一个优化步骤，让我们复制之前的卷积核函数并将其命名为convolution_kernel_v3()：

首先，我们需要预先准备共享内存空间，以便它可以存储输入数据。为了从共享内存中获得滤波操作的好处，我们需要额外的输入数据。为了创建足够的内存空间，我们需要修改内核调用，如下所示：

int shared_mem_size = (2*filter_size+BLOCK_DIM) * (2*filter_size+BLOCK_DIM) * sizeof(float);
convolution_kernel_v3<<<dimGrid, dimBlock, shared_mem_size, 0 >>>(d_output, d_input, d_filter, num_row, num_col, filter_size);

在内核代码中，我们可以声明共享内存空间如下：

extern __shared__ float s_input[];

然后，我们可以将输入数据复制到由线程块计算的共享内存中。首先，让我们声明一些帮助控制内存操作的变量：

int pad_size = filter_size / 2;
int tile_size = BLOCK_DIM + 2 * pad_size;

现在，我们可以按照之前讨论的平铺设计将加载的输入数据复制到共享内存中：

for (int row = 0; row <= tile_size / BLOCK_DIM; row++) {
    for (int col = 0; col <= tile_size / BLOCK_DIM; col++) {
        int idx_row = idx_y + BLOCK_DIM * row - pad_size; 
        // input data index row
        int idx_col = idx_x + BLOCK_DIM * col - pad_size; 
        // input data index column
        int fid_row = threadIdx.y + BLOCK_DIM * row; 
        // filter index row
        int fid_col = threadIdx.x + BLOCK_DIM * col; 
        // filter index column

        if (fid_row >= tile_size || fid_col >= tile_size) continue;

        s_input[tile_size * fid_row + fid_col] = \
            (idx_row >= 0 && idx_row < num_row && idx_col >= 0 
                && idx_col < num_col) ? 
                d_input[num_col * idx_row + idx_col] : 0.f;
    }
}

__syncthreads();

由于输入内存已更改，我们的卷积代码应该更新。我们可以将卷积代码编写如下：

float result = 0.f;
    for (int filter_row = -filter_size / 2; 
         filter_row <= filter_size / 2; ++filter_row)
    {
        for (int filter_col = -filter_size / 2; 
             filter_col <= filter_size / 2; ++filter_col)
        {
            // Find the global position to apply the given filter 
            int image_row = threadIdx.y + pad_size + filter_row;
            int image_col = threadIdx.x + pad_size + filter_col;

            float image_value = s_input[tile_size 
                                        * image_row + image_col]; 
            float filter_value = c_filter[(filter_row 
                                          + filter_size / 2) 
                                          * filter_size 
                                          + filter_col 
                                          + filter_size / 2];

            result += image_value * filter_value;
        }
    }

最后，我们可以使用nvprof来测量性能增益。从结果中，我们可以确认我们的加速速度比原始操作快了大约 35%：

$ nvcc -run -m64 -I/usr/local/cuda/samples/common/inc -gencode arch=compute_70,code=sm_70 -o convolution ./convolution.cu
$ nvprof ./convolution
Processing Time (1) -> GPU: 0.48 ms
Processing Time (2) -> GPU: 0.43 ms
Processing Time (3) -> GPU: 0.30 ms
Processing Time -> Host: 4104.51 ms
... (profiler output) ...
              type Time(%)    Time Calls .    Avg      Min .    Max Name
   GPU activities: 66.85% 2.3007ms     3 766.91us 1.1840us 2.2979ms [CUDA memcpy HtoD]
                   12.85% 442.21us     1 442.21us 442.21us 442.21us convolution_kernel_v1()
                   11.97% 412.00us     1 412.00us 412.00us 412.00us convolution_kernel_v2()
                    8.33% 286.56us     1 286.56us 286.56us 286.56us convolution_kernel_v3()

现在，我们已经了解如何利用加载的数据，以便我们可以重复使用它与其他片上缓存而不是全局内存。我们将在下一节中更详细地讨论这个问题。

获得更多性能

如果滤波器是对称滤波器或可分离滤波器，我们可以将盒式滤波器分解为两个滤波器：水平滤波器和垂直滤波器。使用两个方向滤波器，我们可以在共享内存使用方面进行更多优化：内存空间和内存利用率。如果您想了解更多信息，请查看名为convolutionSeparable的 CUDA 示例，该示例位于3_Imaging/convolutionSeparable目录中。其详细说明也包含在相同目录的doc/convolutionSeparable.pdf中。

前缀和（扫描）

前缀和（扫描）用于从给定的输入数字数组中获得累积数字数组。例如，我们可以按以下方式制作前缀和序列：

输入数字	1	2	3	4	5	6	...
前缀和	1	3	6	10	15	21	...

它与并行减少不同，因为减少只是从给定的输入数据生成总操作输出。另一方面，扫描从每个操作生成输出。解决这个问题的最简单方法是迭代所有输入以生成输出。但是，在 GPU 中这将花费很长时间并且效率低下。因此，温和的方法可以并行化前缀和操作，如下所示：

在这种方法中，我们可以使用多个 CUDA 核心来获得输出。但是，这种方法并不会减少迭代的总次数，因为第一个输入元素应该逐个添加到所有输出中。此外，当数组足够大时，我们无法预测输出结果，因此应该启动多个线程块。这是因为在 CUDA 架构中，并非所有计划的 CUDA 线程都同时启动，并且多个 CUDA 线程会发生冲突。为了避免这种情况，我们需要对数组采用双缓冲区方法，这是另一种低效的方法。以下代码显示了它的实现：

__global__ void
scan_v1_kernel(float *d_output, float *d_input, int length, int offset) {
    int idx = blockDim.x * blockIdx.x + threadIdx.x;

    float element = 0.f;
    for (int offset = 0; offset < length; offset++) {
        if (idx - offset >= 0)
            element += d_input[idx - offset];
    }
    d_output[idx] = element;
}

还有另一种优化方法叫做Blelloch 扫描。该方法通过指数增加和减少步长来生成前缀和输出。该方法的过程如下图所示：

基于步长控制有两个步骤。在增加步长的同时，相应地获得部分总和。然后，在减小步长的同时获得部分总和。每个步骤都有不同的操作模式，但可以根据步长大小来确定。现在，让我们来看一下 Blelloch 扫描的实现并检查更新后的性能。

Blelloch 扫描实现

以下步骤将向您展示如何实现优化的并行扫描算法：

让我们创建一个可以接受输入和输出内存以及它们的大小的内核函数：

__global__ void scan_v2_kernel(float *d_output, float *d_input, int length)
{
    ...
}

然后，我们将创建一个 CUDA 线程索引和一个全局索引来处理输入数据：

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

为了加快迭代速度，我们将使用共享内存。该算法可以生成 CUDA 线程大小的两倍输出，因此我们将额外加载块大小的输入数据到共享内存中：

extern __shared__ float s_buffer[];
s_buffer[threadIdx.x] = d_input[idx];
s_buffer[threadIdx.x + BLOCK_DIM] = d_input[idx + BLOCK_DIM];

在开始迭代之前，我们将声明偏移变量，该变量计算左操作数和右操作数之间的差距：

int offset = 1;

然后，我们将添加输入数据，直到偏移量大于输入的长度为止：

while (offset < length)
{
    __syncthreads();
    int idx_a = offset * (2 * tid + 1) - 1;
    int idx_b = offset * (2 * tid + 2) - 1;
    if (idx_a >= 0 && idx_b < 2 * BLOCK_DIM) {
        s_buffer[idx_b] += s_buffer[idx_a];
    }
    offset <<= 1;
}

之后，我们将通过减小减少大小来再次迭代两次：

offset >>= 1;
while (offset > 0) {
    __syncthreads();
    int idx_a = offset * (2 * tid + 2) - 1;
    int idx_b = offset * (2 * tid + 3) - 1;
    if (idx_a >= 0 && idx_b < 2 * BLOCK_DIM) {
        s_buffer[idx_b] += s_buffer[idx_a];
    }
    offset >>= 1;
}
__syncthreads();

最后，我们将使用内核函数将输出值存储在全局内存中：

d_output[idx] = s_buffer[tid];
d_output[idx + BLOCK_DIM] = s_buffer[tid + BLOCK_DIM];

现在，我们可以调用这个扫描内核函数如下：

void scan_v2(float *d_output, float *d_input, int length)
{
    dim3 dimBlock(BLOCK_DIM);
    dim3 dimGrid(1);
    scan_v2_kernel<<<dimGrid, dimBlock, 
                     sizeof(float) * BLOCK_DIM * 2>>>
                  (d_output, d_input, length);
    cudaDeviceSynchronize();
}

您还可以使用相同的函数接口编写一个朴素扫描版本。现在，让我们回顾一下我们的新版本有多快，以及我们是否可以利用其他优化机会。

以下代码显示了朴素扫描和 Blelloch 扫描性能的分析结果：

$ nvcc -m64 -std=c++11 -I/usr/local/cuda/samples/common/inc -gencode arch=compute_70,code=sm_70 -L/usr/local/cuda/lib -o scan ./scan.cu ./scan_v1.cu ./scan_v2.cu
$ nvprof ./scan
            Type Time(%)     Time Calls      Avg      Min      Max Name
 GPU activities:  68.96% 22.751us     1 22.751us 22.751us 22.751us scan_v1_kernel(float*, float*, int)
 12.71% 4.1920us    1 4.1920us 4.1920us 4.1920us scan_v2_kernel(float*, float*, int)

正如你所看到的，由于减少了开销，Blolloch 扫描比朴素扫描算法快了大约五倍。我们还可以通过比较不同实现的输出来验证操作结果：

input         :: -0.4508 -0.0210 -0.4774  0.2750 ... 0.0398 0.4869
result[cpu]   :: -0.4508 -0.4718 -0.9492 -0.6742 ... 0.3091 0.7960
result[gpu_v1]:: -0.4508 -0.4718 -0.9492 -0.6742 ... 0.3091 0.7960
SUCCESS!!
result[cpu]   :: -0.4508 -0.4718 -0.9492 -0.6742 ... 0.3091 0.7960
result[gpu_v2]:: -0.4508 -0.4718 -0.9492 -0.6742 ... 0.3091 0.7960
SUCCESS!!

到目前为止，我们已经介绍了如何设计和实现优化的单个块大小的并行前缀和操作。要在输入数据上使用前缀和操作，需要比块大小更多的数据，我们需要基于我们的块级减少代码构建一个块级前缀和操作。我们将在下一节详细讨论这个问题。

构建全局大小扫描

我们实现的前缀和操作在单个线程块内工作。由于第一步有两个输入，而我们在一个线程块中最多可以有 1,024 个 CUDA 线程，因此最大可用大小为 2,048。在不考虑其他线程块操作的情况下，线程块进行上扫描和下扫描。

然而，如果我们执行一个分块扫描操作，这个操作可以被扩大。为了做到这一点，你需要额外的步骤来收集最后一个前缀和的结果，扫描它们，并将每个线程块的结果与每个块的块级扫描值相加。这个过程可以按照以下方式实现：

追求更好的性能

我们的实现代码执行了最佳操作。然而，我们可以通过减少共享内存的银行冲突来进一步优化。在我们的实现中，CUDA 线程在某些点上访问相同的内存银行。NVIDIA 的 GPU Gem3 在第三十九章，使用 CUDA 进行并行前缀和（扫描）（developer.nvidia.com/gpugems/GPUGems3/gpugems3_ch39.html）中介绍了前缀和（扫描），并在39.2.3 避免银行冲突中指出了这个问题。你可以将解决方案调整到我们的实现，但如果这样做，你应该将NUM_BANKS更新为32，LOG_NUM_BANKS更新为5。现在，CUDA 架构有 32 个共享内存银行。

并行前缀和操作的其他应用

G.E. Blelloch 博士在 1993 年发表了一篇关于他的前缀和算法的文章前缀和及其应用（www.cs.cmu.edu/~guyb/papers/Ble93.pdf）。通过阅读他的文章，你可以了解更多关于并行前缀和算法及其应用。这些应用包括压缩、分割、分段扫描、快速排序、基数排序和归并排序。

Ahmed Sallm 博士的视频讲座，使用 CUDA 进行并行处理简介-第 4 讲第 2\3 部分（youtu.be/y2HzWKTqo3E），对此提供了很好的介绍。它提供了关于前缀和算法如何用于裁剪图形和构建稀疏矩阵的概念介绍。他还提供了关于如何使用排序算法的说明。

压缩和分割

之前，我们介绍了如何并行化顺序前缀和算法，并讨论了它如何用于其他应用。现在，让我们来介绍其中一些应用：压缩和分割。压缩操作是一种可以从数组中整合满足给定条件的值的算法。另一方面，分割操作是一种将值分配到指定位置的算法。一般来说，这些算法是顺序工作的。然而，我们将看到并行前缀和操作如何改进它的功能。

压缩操作用于将满足特定条件的特定数据收集到一个数组中。例如，如果我们想要对数组中的正元素使用压缩操作，那么操作如下：

在并行编程中，我们有一种不同的方法，可以利用并行前缀和操作使用多个核心。首先，我们标记数据以检查它是否满足条件（即谓词），然后进行前缀和操作。前缀和的输出将是标记值的索引，因此我们可以通过复制它们来获得收集的数组。下面的图表显示了压缩操作的一个示例：

由于所有这些任务都可以并行完成，我们可以在四个步骤中获得收集的数组。

另一方面，拆分意味着将数据分发到多个不同的位置。一般来说，我们会从最初的位置分发数据。下面的图表显示了它的操作示例：

这个例子显示了收集的数组元素是如何分布在它们原来的位置的。我们也可以使用前缀和并行地做到这一点。首先，我们参考谓词数组并进行前缀和操作。由于输出是每个元素的地址，我们可以很容易地分配它们。下面的图表显示了如何进行这个操作：

现在，让我们实现这个并讨论它们的性能限制和应用。

实现压缩

压缩操作是一个谓词、扫描、寻址和收集的序列。在这个实现中，我们将从一个随机生成的数字数组中构建一个正数数组。初始版本只能承受单个线程块操作，因为我们只会使用一个块大小的前缀和操作。然而，我们可以了解前缀和如何对其他应用有用，并将这个操作扩展到更大的数组，使用扩展的前缀和操作。

为了实现压缩操作，我们将编写几个核函数，可以为每个步骤执行所需的操作，并调用最后那些：

让我们编写一个核函数，通过检查每个元素的值是否大于零来生成一个谓词数组：

__global__ void
predicate_kernel(float *d_predicates, float *d_input, int length)
{
    int idx = blockDim.x * blockIdx.x + threadIdx.x;

    if (idx >= length) return;

    d_predicates[idx] = d_input[idx] > FLT_ZERO;
}

然后，我们必须对谓词数组执行前缀和操作。我们将在这里重用之前的实现。之后，我们可以编写一个可以检测扫描数组的地址并将目标元素收集为输出的核函数：

__global__ void
pack_kernel(float *d_output, float *d_input, float *d_predicates, float *d_scanned, int length)
{
    int idx = blockDim.x * blockIdx.x + threadIdx.x;

    if (idx >= length) return;

    if (d_predicates[idx] != 0.f)
    {
        // addressing
        int address = d_scanned[idx] - 1;

        // gather
        d_output[address] = d_input[idx];
    }
}

现在，让我们一起调用它们来进行压缩操作：

// predicates
predicate_kernel<<< GRID_DIM, BLOCK_DIM >>>(d_predicates, d_input, length);
// scan
scan_v2(d_scanned, d_predicates, length);
// addressing & gather (pack)
pack_kernel<<< GRID_DIM, BLOCK_DIM >>>(d_output, d_input, d_predicates, d_scanned, length);

现在，我们有了一个从随机生成的数组中收集到的正数数组：

$ nvcc -run -m64 -std=c++11 -I/usr/local/cuda/samples/common/inc -gencode arch=compute_70,code=sm_70 -L/usr/local/cuda/lib -o pack_n_split ./pack_n_split.cu
input    :: -0.4508 -0.0210 -0.4774  0.2750 .... 0.0398  0.4869
pack[cpu]::  0.2750  0.3169  0.1248  0.4241 .... 0.3957  0.2958
pack[gpu]::  0.2750  0.3169  0.1248  0.4241 .... 0.3957  0.2958
SUCCESS!!

通过使用并行前缀和操作，我们可以很容易地并行实现压缩操作。我们的实现从给定数组中压缩正值，但我们可以将其切换到其他条件并且应用压缩操作而不会有困难。现在，让我们来讨论如何将这些压缩元素分发到原始数组。

实现拆分

拆分操作是一个谓词、扫描、地址和拆分的序列。在这个实现中，我们将重用在前一节中创建的地址数组。因此，我们可以跳过之前的步骤，只需从地址数组中实现拆分操作：

让我们编写拆分核函数，如下所示：

__global__ void
split_kernel(float *d_output, float *d_input, float *d_predicates, float *d_scanned, int length)
{
    int idx = blockDim.x * blockIdx.x + threadIdx.x;

    if (idx >= length) return;

    if (d_predicates[idx] != 0.f)
    {
        // address
        int address = d_scanned[idx] - 1;

        // split
        d_output[idx] = d_input[address];
    }
}

现在，我们可以调用核函数，如下所示：

cudaMemcpy(d_input, d_output, sizeof(float) * length, cudaMemcpyDeviceToDevice);
    cudaMemset(d_output, 0, sizeof(float) * length);
    split_kernel<<<GRID_DIM, BLOCK_DIM>>>(d_output, d_input, d_predicates, d_scanned, length);

由于我们将使用前一步骤的扫描输出，我们将把它复制到输入并清除原始数组。总的来说，我们可以使用 CUDA 进行并行压缩和拆分。这是我们实现的输出。您可以确认它按预期运行：

$ nvcc -run -m64 -std=c++11 -I/usr/local/cuda/samples/common/inc -gencode arch=compute_70,code=sm_70 -L/usr/local/cuda/lib -o pack_n_split ./pack_n_split.cu
input    :: -0.4508 -0.0210 -0.4774  0.2750 .... 0.0398  0.4869
pack[cpu]::  0.2750  0.3169  0.1248  0.4241 .... 0.3957  0.2958
pack[gpu]::  0.2750  0.3169  0.1248  0.4241 .... 0.3957  0.2958
SUCCESS!!
split[gpu]   0.0000  0.0000  0.0000  0.2750 .... 0.0398  0.4869
SUCCESS!!

在我们的实现中，我们为正值生成了一个压缩数组和一个拆分数组。由于并行前缀和，我们也可以并行地做到这一点。我们版本的一个主要限制是，它只支持少于 2,048 个元素，因为我们的实现是基于之前的并行前缀和实现的。

N-body

任何 N 体模拟都是一个在物理力的影响下演化的动力学系统的模拟。随着物体不断相互作用，进行数值近似。N 体模拟在物理学和天文学中被广泛使用，例如，科学家可以了解宇宙中粒子的动态。N 体模拟也在许多其他领域中使用，包括计算流体动力学，以便理解湍流流体流动模拟。

解决 N 体模拟的一个相对简单的方法是利用O(N²)复杂度的蛮力技术。这种方法在本质上是尴尬地并行的。在算法规模上有各种优化可以减少计算复杂度。可以用来确定近距离相互作用中的力，而不是将所有对应用于整个模拟。即使在这种情况下，为 CUDA 解决力量创建一个内核也是非常有用的，因为它还将提高远场组件的性能。加速一个组件将卸载其他组件的工作，因此整个应用程序都会从加速一个内核中受益。

在 GPU 上实现 N 体模拟

该算法基本上是一个计算力f[ij]的所有对算法，对于一个 NN 网格。一个物体i上的总力/加速度F[i]是该行中所有条目的总和。从并行性的角度来看，这是一个尴尬地并行的任务，复杂度为O(N²)。

从性能角度来看，该应用程序受到内存限制，并且会受到内存带宽的限制。好的一点是，许多数据可以被重复使用并存储在高带宽和低延迟的内存中，比如共享内存。在共享内存中重复使用和存储数据可以减少对全局内存的负载，从而有助于达到峰值计算性能。

以下图表显示了我们将使用的策略：

我们不再从全局内存中反复加载内存，而是利用平铺。我们已经在之前的章节中演示了在矩阵乘法中使用平铺，并在图像应用中使用了它。前面的图表显示了每一行都是并行评估的。平铺大小由可以存储在共享内存中而不影响内核占用率的最大元素数量定义。每个块将数据加载到共享内存中，然后执行同步。一旦数据加载到共享内存中，就在每个块中进行力/加速度计算。可以看到，即使单独的行是并行计算的，为了实现最佳的数据重用，每行中的相互作用是顺序进行的。

N 体模拟实现概述

让我们以伪代码格式回顾一下这个实现，然后解释它的逻辑。在这个例子中，我们使用引力势来说明所有对 N 体模拟中的基本计算形式。实现的代码可以在07_parallel_programming_pattern/05_n-body中找到。按照以下步骤开始：

用随机变量初始化 n 空间：

data[i] = 2.0f * (rand() / max) - 1.0f

在一个中间共享内存空间中声明和存储数据，以便有效地重复使用。同步以确保块内的所有线程都能看到共享内存中的更新值：

for (int tile = 0; tile < gridDim.x; tile++) {
... 
__shared__ float3 shared_position[blockDim.x];
float4 temp_position = p[tile * blockDim.x + threadIdx.x];
shared_position[threadIdx.x] = make_float3(temp_position.x, temp_position.y, temp_position.z);
__syncthreads();
...
}

通过迭代每个块来计算力：

for (int j = 0; j < BLOCK_SIZE; j++) {
    //Calculate Force
    __syncthreads();
}

最后，使用以下命令将应用程序编译为nvcc编译器：

$nvcc -run --gpu-architecture=sm_70 -o n-body n_body.cu

正如你所看到的，实现 N 体模拟是一个尴尬地并行的任务，而且非常简单。虽然我们在这里实现了基本版本的代码，但存在各种算法变体。你可以利用这个版本作为一个模板，根据对算法的更改进行改进。

直方图计算

在一个尴尬的并行作业中，理想情况下，您会将计算分配给每个线程，这些线程在独立数据上工作，从而不会发生数据竞争。到目前为止，您可能已经意识到有些模式不适合这个类别。其中一种模式是当我们计算直方图时。直方图模式显示了数据项的频率，例如，我们在每个章节中使用 CUDA 这个词的次数

章节，本章中每个字母出现的次数等。直方图采用以下形式：

在这一部分，我们将利用原子操作来串行访问数据，以便获得正确的结果。

编译和执行步骤

直方图提供了关于手头数据集的重要特征，以及有用的见解。例如，在整个图像中，只有少数区域可能存在感兴趣的区域。有时创建直方图用于找出图像中可能存在感兴趣区域的位置。在这个例子中，我们将使用在整个图像中将图像分成块来计算直方图。让我们开始吧：

准备您的 GPU 应用程序。此代码可以在07_parallel_programming_pattern/08_histogram中找到。
使用以下命令将您的应用程序编译为nvcc编译器：

$ nvcc -c scrImagePgmPpmPackage.cpp 
$ nvcc -c image_histogram.cu
$ nvcc -run -o image_histogram image_histogram.o scrImagePgmPpmPackage.o

scrImagePgmPpmPackage.cpp文件提供了我们可以用来读取和写入.pgm扩展名图像的源代码。直方图计算代码可以在image_histogram.cu中找到。

理解并行直方图

诸如直方图之类的模式需要原子操作，这意味着以串行方式更新特定地址的值，以消除多个线程之间的争用，从而更新相同的地址。这需要多个线程之间的协调。在这个七步过程中，您可能已经注意到我们使用了私有化。私有化是一种利用低延迟内存（如共享内存）来减少吞吐量和降低延迟的技术，如下图所示：

基本上，我们不是在全局内存上使用原子操作，而是在共享内存上使用原子操作。原因现在应该对您来说是相当明显的。与在共享内存/ L1 缓存上执行相同操作相比，在全局内存上执行原子操作的成本更高。从 Maxwell 架构开始，原子操作得到了硬件支持。私有化的共享内存实现应该从 Maxwell 架构开始为您提供 2 倍的性能。但是，请注意，原子操作仅限于特定的函数和数据大小。

使用 CUDA 原子函数计算直方图

主要地，我们将利用共享内存上的atomicAdd()操作来计算共享内存中每个块的直方图。按照以下步骤在内核中计算直方图：

为每个块分配与每个块的直方图大小相等的共享内存。由于这是一个 char 图像，我们期望元素在 0-255 的范围内：

__shared__ unsigned int histo_private[256];

将每个块的共享内存数组初始化为0：

if(localId <256)
    histo_private[localId] = 0;

同步这一点，以确保块内的所有线程看到初始化的数组：

__syncthreads();

从全局/纹理内存中读取图像的数据：

unsigned char imageData = tex2D<unsigned char>(texObj,(float)(tidX),(float)(tidY));

在共享内存上进行atomicAdd()操作：

atomicAdd(&(histo_private[imageData]), 1);

在写入全局内存之前，在块之间进行同步：

__syncthreads();

将每个块的直方图写入全局内存：

if(localId <256)
    imageHistogram[histStartIndex+localId] = histo_private[localId];

现在，我们已经完成了在 GPU 上实现直方图计算。

总之，使用共享原子内存很容易实现直方图。由于硬件对共享原子内存的本机支持，这种方法可以在 Maxwell 架构之后的显卡上获得高性能。

使用动态并行性在 CUDA 中进行快速排序

作为任何应用程序的基本构建块的关键算法之一是排序。有许多可用的排序算法已经得到了广泛的研究。最坏时间复杂度、最佳时间复杂度、输入数据特征（数据几乎排序好还是随机的？是键值对吗？是整数还是浮点数？）、原地或非原地内存需求等等，这些都定义了哪种算法适用于哪种应用。一些排序算法属于分治算法的范畴。这些算法适合并行处理，并适用于 GPU 等可以将要排序的数据分割进行排序的架构。其中一个这样的算法是快速排序。正如我们之前所述，快速排序属于分治范畴。它是一个三步方法，如下：

从需要排序的数组中选择一个元素。这个元素作为枢轴元素。
第二步是分区，确定所有元素的位置。所有小于枢轴的元素都移到左边，所有大于或等于枢轴的元素都移到枢轴元素的右边。这一步也被称为分区。
递归地执行步骤 1 和 2，直到所有子数组都被排序。

快速排序的最坏情况复杂度是 O(n²)，这与其他排序过程的最坏情况复杂度为 O(nlogn)相比可能不太理想（例如归并排序和堆排序）。然而，实际上快速排序被认为是有效的。枢轴元素的选择可以经过考虑，有时也可以随机选择，以使最坏情况复杂度几乎不会发生。此外，与其他排序算法相比，快速排序的内存负载和需求较少，例如归并排序需要额外的存储空间。更实际的快速排序实现使用随机化版本。随机化版本的期望时间复杂度为 O(nlogn)。最坏情况复杂度在随机化版本中也是可能的，但它不会发生在特定模式（例如排序好的数组）上，随机化快速排序在实践中表现良好。

虽然我们可以写一整章关于排序算法的特性，但我们计划只覆盖 CUDA 的特性，这将帮助您在 GPU 上高效实现快速排序。在本节中，我们将使用从 CUDA 6.0 和 GPU 架构 3.5 开始引入的动态并行性。

现在，让我们回顾一下动态并行性是如何对排序算法做出贡献的。

快速排序和 CUDA 动态并行性

快速排序算法要求递归地启动内核。到目前为止，我们所见过的算法是通过 CPU 一次调用内核。内核执行完毕后，我们返回到 CPU 线程，然后重新启动它。这样做会导致将控制权交还给 CPU，并且可能导致 CPU 和 GPU 之间的数据传输，这是一项昂贵的操作。以前在 GPU 上高效实现需要递归等特性的算法（如快速排序）曾经非常困难。从 GPU 架构 3.5 和 CUDA 5.0 开始，引入了一个名为动态并行性的新特性。

动态并行性允许内核内的线程在不将控制权返回给 CPU 的情况下从 GPU 上启动新的内核。动态一词来自于它基于运行时数据的动态性。多个内核可以同时由线程启动。以下图表简化了这个解释：

如果我们将这个概念转化为快速排序的执行方式，它会看起来像这样：

深度 0 是来自 CPU 的调用。对于每个子数组，我们启动两个内核：一个用于左数组，一个用于右数组。递归在达到内核的最大深度或元素数量小于 32（即 warp 大小）后停止。为了使内核的启动在非零流中是异步的，以便子数组内核可以独立启动，我们需要在每次内核启动之前创建一个流：

cudaStream_t s;
cudaStreamCreateWithFlags( &s, cudaStreamNonBlocking );
cdp_simple_quicksort<<< 1, 1, 0, s >>>(data, left, nright, depth+1);
cudaStreamDestroy( s );

这是一个非常重要的步骤，否则内核的启动可能会被序列化。有关流的更多细节，请参考多 GPU 内核。

CUDA 的 Quicksort

对于我们的 Quicksort 实现，我们将利用动态并行性来递归启动 GPU 内核。实现 Quicksort 的主要步骤如下：

CPU 启动第一个内核：内核以一个块和一个线程启动。左元素是数组的开始，右元素是数组的最后一个元素（基本上是整个数组）：

int main(int argc, char **argv)
{ ...
    cdp_simple_quicksort<<< 1, 1 >>>(data, left, right, 0);
}

限制检查：在从内核内部启动内核之前检查两个条件。首先，检查我们是否已经达到硬件允许的最大深度限制。其次，我们需要检查子数组中要排序的元素数量是否小于 warp 大小（32）。如果其中一个条件为真，那么我们必须按顺序执行选择排序，而不是启动一个新的内核：

__global__ void cdp_simple_quicksort( unsigned int *data, int left, int right, int depth )
{ ...

if( depth >= MAX_DEPTH || right-left <= INSERTION_SORT )
 {
     selection_sort( data, left, right );
     return;
 }

分区：如果满足前面的条件，那么将数组分成两个子数组，并启动两个新的内核，一个用于左数组，另一个用于右数组。如果你仔细看下面的代码，你会发现我们是从内核内部启动内核的：

__global__ void cdp_simple_quicksort( unsigned int *data, int left, int right, int depth ) {
...
while(lptr <= rptr)
 {
     // Move the left pointer as long as the 
     // pointed element is smaller than the pivot.
     // Move the right pointer as long as the 
     // pointed element is larger than the pivot.
     // If the swap points are valid, do the swap!

     // Launch a new block to sort the left part.
     if(left < (rptr-data))
     { // Create a new stream for the eft sub array
        cdp_simple_quicksort<<< 1, 1, 0, s 
                            >>>(data, left, nright, depth+1);
     }
    // Launch a new block to sort the right part.
    if((lptr-data) < right)
     {//Create stream for the right sub array
         cdp_simple_quicksort<<< 1, 1, 0, s1 
                             >>>(data, nleft, right, depth+1);
     }
 }

执行代码：实现的代码可以在07_parallel_programming_pattern/06_quicksort中找到。使用以下命令使用nvcc编译您的应用程序：

$nvcc -o quick_sort --gpu-architecture=sm_70 -rdc=true quick_sort.cu

如你所见，我们在编译中添加了两个标志：

-- gpu-architecture=sm_70：这个标志告诉nvcc为 Volta GPU 编译和生成二进制/ptx。如果你没有特别添加这个标志，编译器会尝试从sm_20（即 Fermi 代）兼容的代码编译，直到新架构sm_70（即 Volta）。由于旧一代的卡不支持动态并行性，编译将失败。
-rdc=true：这是一个关键参数，它在 GPU 上启用动态并行性。

动态并行性指南和约束

虽然动态并行性为我们提供了在 GPU 上移植 Quicksort 等算法的机会，但需要遵循一些基本规则和指南。

编程模型规则：基本上，所有 CUDA 编程模型规则都适用：

内核启动是每个线程异步的。
同步只允许在块内进行。
创建的流在一个块内共享。
事件可用于创建流间依赖关系。

内存一致性规则：

子内核在启动时看到父内核的状态。
父内核只能在同步后看到子内核所做的更改。
本地和共享内存通常是私有的，父内核无法传递或访问。

指南：

重要的是要理解，每次内核启动都会增加延迟。从另一个内核内部启动内核的延迟随着新架构的推出逐渐减少。
虽然启动吞吐量比主机高一个数量级，但最大深度可以设置限制。最新一代卡允许的最大深度是 24。
从内核内部执行cudaDeviceSynchronize()是一个非常昂贵的操作，应尽量避免。
在全局内存上预先分配了额外的内存，以便在启动之前存储内核。
如果内核失败，错误只能从主机上看到。因此，建议您使用-lineinfo标志以及cuda-memcheck来定位错误的位置。

基数排序

另一个非常流行的排序算法是基数排序，因为它在顺序机器上非常快。基数排序的基本策略是每个元素都按位排序。让我们看一个简单的例子来解释基数排序涉及的步骤：

假设要排序的元素如下：

值	7	14	4	1

这些数字的等效二进制值如下：

位	0111	1110	0100	0001

第一步是根据第 0 位进行排序。这些数字的第 0 位如下：

0 位	1	0	0	1

根据第 o 位排序基本上意味着所有的零都在左边。所有的 1 都在右边，同时保持元素的顺序：

第 0 位上的排序值	14	4	7	1
根据第 0 位排序的位	1110	0100	0111	0001

第 0 位完成后，我们继续到第一位。根据第一位排序后的结果如下：

第一位上的排序值	4	14	7	1
根据第一位排序的位	0100	1110	0111	0001

然后，我们继续到下一个更高的位，直到所有的位都结束。最终结果如下：

所有位上的排序值	1	4	7	1
根据所有位排序的位	0001	0100	0111	1110

正如您所看到的，在这个例子中我们设置的上限是 4 位。对于更大的数字，比如整数，这将持续到 32 位，因为整数是 32 位的。

现在我们已经了解了这个算法，让我们看看如何在 GPU 中实现它。与本章中的其他部分相比，我们将采取两种方法来展示 CUDA 生态系统，以便我们可以实现/使用基数排序。

选项 1：我们将使用翘曲级别来对 32 个元素进行基数排序。这样做的原因是我们希望利用基数排序来向您介绍翘曲级别原语。

选项 2：我们将使用 CUDA 工具包的一部分 Thrust 库。它实现了通用基数排序。最好的实现是重用。由于 Thrust 已经提供了最好的基数排序实现之一，我们将使用它。

两种方法

为了方便您的理解，让我们从示例代码开始。在这个例子中，我们将使用翘曲级别原语和 Thrust 库来实现/使用基数排序。示例代码可以在07_parallel_programming_pattern/07_radixsort中找到。

使用以下命令使用nvcc编译器编译您的应用程序：

翘曲级别原语版本：

$ nvcc -run -o radix_warp_sort radix_warp_sort.cu

Thrust 库版本：

$ nvcc -run -o radix_thrust_sort thrust_radix_sort.cu

这两个例子展示了 GPU 给出的排序输出。现在，让我们详细了解这些操作是如何实现的。

方法 1 - 翘曲级别原语

让我们看看 CUDA 翘曲级别原语是如何在代码中实现我们的算法的：

首先，将数据从全局内存加载到共享内存中：

__shared__ unsigned int s_data[WARP_SIZE*2];

内存的大小等于翘曲大小，*2，以便它可以实现乒乓缓冲区。

从低位到高位循环：

for (int i = MIN_BIT_POS; i <= MAX_BIT_POS; i++){ ... }

获取当前的掩码：

unsigned int bit  = data&bit_mask;

获取 1 和 0 的数量（直方图）：

unsigned int active = __activemask();
unsigned int ones = __ballot_sync(active,bit);
unsigned int zeroes = ~ones;

获取当前位数为零（0）的线程的位置（前缀和）。
获取当前位数为一（1）的线程的位置（前缀和）：

if (!bit) // threads with a zero bit
 // get my position in ping-pong buffer
 pos = __popc(zeroes&thread_mask);
 else // threads with a one bit
 // get my position in ping-pong buffer
 pos = __popc(zeroes)+__popc(ones&thread_mask);

将数据存储在乒乓共享缓冲区内存中：

 s_data[pos-1+offset] = data;

重复步骤 2-6，直到达到上限位。
从共享内存中将最终结果存储到全局内存中：

d_data[threadIdx.x] = s_data[threadIdx.x+offset];

也许对于您来说，直方图和前缀和突然出现可能不太清楚。让我们详细讨论这个实现，以便我们可以理解如何使用翘曲级别原语来实现相同的功能。

在本节的开头，我们描述了如何使用示例进行排序。然而，我们没有涵盖的是如何找出需要交换的元素的位置。基数排序可以使用基本原语（如直方图和前缀和）来实现，因此可以很容易地在 GPU 上实现。

让我们重新审视我们看过的示例，并收集其细节，包括直方图和前缀和的步骤。以下表格显示了在每个位上迭代进行的各种计算：

值	7	14	4	1
二进制	0111	1110	0100	0001
位 0	1	0	0	1
直方图前缀和	2	0	2	2
偏移	0	0	1	1
新索引（前缀和和偏移）	2	0	1	3

让我们解释前面表格中显示的每一项计算，如下所示：

首先，我们为第 0 位位置的元素构建直方图，包括具有 0 和 1 的元素的数量：

直方图：零位（2 个值），一位（2 个值）

然后，我们对这些值进行排他性前缀和。前缀和可以定义为所有先前值的总和。在我们的情况下，我们分别对 0 位和 1 位进行这样的操作。
最后，我们根据前缀和的值移动元素。

我们用来找到直方图和前缀和的 warp 级原语分别是__ballot_sync()和__popc()。

__ballot_sync() API 评估 warp 的所有活动线程的谓词，并返回一个整数，其第 N 位设置为 1，当且仅当谓词对于 warp 的第 N 个线程求值为非零时。__popc()用于计算整数的数量，被设置为 1。

在 CUDA 编程模型中，我们已经看到最小的执行单元是一个 warp（32 个线程）。CUDA 提供了各种 warp 级原语，可以进行细粒度控制，在许多应用中可以实现更好的性能。我们在上一节介绍了一个这样的原语__ballot_sync()。其他重要的 warp 级原语包括shuffle指令，用于特定的 warp 级归约。shuffle指令已经在本书中介绍过。如果您已经达到了 CUDA 的忍者程序员水平，那么我们建议您查看 CUDA API 指南，以了解更多这些 warp 级原语。

这完成了使用 warp 级原语描述基数排序。现在，让我们看看基于 Thrust 库的实现。

方法 2 - 基于 Thrust 的基数排序

基于 Thrust 的基数排序是基数排序的通用实现，对于不同类型的数据（如整数、浮点数或键值对）都能很好地工作。我们想再次强调排序是一个经过深入研究的算法，因此有其并行实现。因此，我们建议在自己实现之前重用现有的库。

使用 Thrust 进行基数排序的步骤如下：

导入相关的头文件（Thrust 是一个仅包含头文件的库，类似于 STL）：

#include <thrust/device_vector.h>
#include <thrust/sort.h>

声明并初始化设备向量：

//declare a device vector of size N
thrust::device_vector<int> keys(N);
//Generate a random number generator engine
thrust::default_random_engine r(12);
//create a distribution engine which will create integer values
thrust::uniform_int_distribution<int> d(10, 99);
//Fill the array with randon values
for(size_t i = 0; i < v.size(); i++)
    v[i] = d(r);

对初始化的设备向量进行排序：

thrust::sort(keys.begin(), keys.end());

使用这个库提供了一种更简单和更健壮的方法。Thrust 提供了不同类型的排序方法，包括整数和浮点数的基数排序。或者，您可以创建一个自定义比较器来进行自定义排序，例如按照偶数后面是奇数的顺序排序，按降序排序等等。如果您想了解更多关于基于 Thrust 的排序示例，建议您查看 CUDA 提供的示例示例。

现在，我们已经看过了在 GPU 上实现基数排序的两种方法。

总结

在本章中，我们看了 CUDA 中常用算法和模式的实现。这些算法和模式是常见的。我们涵盖了矩阵乘法和卷积滤波中的基本优化技术。然后，我们扩展了讨论，介绍了如何通过使用前缀和、N 体、直方图和排序来并行化问题。为此，我们使用了专门的 GPU 知识、库和较低级别的原语。

我们所涵盖的许多算法都是在 CUDA 库中实现的。例如，矩阵乘法在 cuBLAS 库中，而卷积在 CUDNN 库中。此外，我们还涵盖了基数排序实现中的两种方法：使用 Thrust 库或 warp 级原语进行直方图计算。

现在您已经看到了这些模式如何在常用库中实现，下一个合乎逻辑的步骤是看看我们如何可以使用这些库。这就是我们将在下一章中要做的事情。

第八章：使用库和其他语言进行编程

本章还涵盖了其他 GPU 编程方法——使用 GPU 加速库和其他语言进行编程。使用 GPU 加速库进行编程使我们能够开发具有优化内核的应用程序。此外，我们可以使用其他编程语言开发了解 CUDA 加速的 CUDA 软件。这两种方式都提高了可编程性和生产力。此外，我们不必花时间优化已经优化的常见操作。

CUDA 工具包提供了许多线性代数、图像和信号处理以及随机处理的 GPU 加速库。它们包括 cuBLAS（基本线性代数子程序）、cuFFT（快速傅里叶变换）、cuRAND（随机数生成）、NPP（图像和信号处理）、cuSPARSE（稀疏线性代数）、nvGRAPH（图分析）、cuSolver（GPU 中的 LAPACK）、Thrust（CUDA 中的 STL）等。我们还可以使用 OpenCV 库编写 GPU 加速程序。本章将涵盖其中一些库。

我们还可以在 R、MATLAB、Octave 和 Python 中使用 GPU 加速。如今，Python 集成很受欢迎且功能强大，因为 GPU 可以加速许多机器学习和数据科学任务。我们也将这些语言作为入门级内容进行介绍。

本章将涵盖以下主题：

使用 cuBLAS 进行线性代数运算
使用 cuBLAS 进行混合精度运算
用于并行随机数生成的 cuRAND
用于 GPU 的快速傅里叶变换的 cuFFT
使用 GPU 进行图像和信号处理的 NPP
在 OpenCV 中编写 GPU 加速代码
编写与 CUDA 配合使用的 Python 代码
Octave 和 R 中的零编码加速 NVBLAS
MATLAB 中的 CUDA 加速

使用 cuBLAS 进行线性代数运算

cuBLAS 库是 GPU 优化的基本线性代数子程序（BLAS）的标准实现。使用其 API，程序员可以将计算密集型代码优化为单个 GPU 或多个 GPU。cuBLAS 有三个级别。级别 1 执行矢量-矢量运算，级别 2 执行矩阵-矢量运算，级别 3 执行矩阵-矩阵运算。

涵盖每个级别超出了本书的范围。我们只关注如何使用 cuBLAS API 并将其性能扩展到多个 GPU。具体来说，本文将涵盖单精度浮点矩阵乘法（SGEMM）运算——一个三级运算。

cuBLAS 库是 CUDA 工具包的一部分，因此您可以在不进行额外安装的情况下使用 cuBLAS。此外，您可以使用cc或cpp文件扩展名，而不是.cu，因为您不需要使用 CUDA 特定的内置关键字，如__global__或threadIdx。以下代码片段显示了 cuBLAS 函数（cubalsSgemm）的基本应用：

cublasHandle_t handle;
cublasCreate(&handle);

.. { data operation } ..

cublasSgemm(...);

.. { data operation } ..

cublasDestroy(handle);

正如你所看到的，cuBLAS API 使用cublasHandle_t类型的句柄。

cuBLAS SGEMM 运算

GEMM 运算可以用以下方程表示：

其中alpha和beta是标量，A、B和C是以列为主的矩阵。这与以下框中 cuBLAS 函数接口相匹配：

cublasStatus_t cublasSgemm(cublasHandle_t handle,
                          cublasOperation_t transa, 
                           cublasOperation_t transb,
                           int m, int n, int k,
                           const float *alpha,
                           const float *A, int lda,
                           const float *B, int ldb,
                           const float *beta,
                           float *C, int ldc);

在使用这个 GEMM 函数之前，让我们看一下参数的细节：

transa和transb：cuBLAS 函数的指令，用于确定是否应该对矩阵A和B进行转置操作。
m，n和k：矩阵的维度大小。
alpha和beta：确定如何配置源的输出值的参数。
*A，*B和*C：矩阵数据的线性缓冲区。
lda：矩阵A的主列维度。cuBLAS 将矩阵元素与此值对齐。
ldb：矩阵B的主列维度。cuBLAS 将矩阵元素与此值对齐。

要在主机和设备之间传输数据，您可以使用 cuBLAS 的cublasSetMatrix()和cublasGetMatrix()辅助函数。它们是cudaMemcpy()的包装函数，但具有矩阵的维度信息，因此它们有助于增强代码的可读性；当然，您也可以简单地使用cudaMemcpy()。

让我们实现一个应用程序，使用 cuBLAS SGEMM 函数进行 GEMM 操作。我们将包括cublas_v2.h以使用更新的 cuBLAS API。为了方便起见，我们将使用getMatrix()函数从给定的维度获取一个随机生成的矩阵，并使用printMatrix()函数打印矩阵元素。代码已在给定的示例代码中实现。在主函数中，我们将从给定的M、N和K初始化三个矩阵—A、B和C。然后我们将计算cublasSgemm()如下：

cublasHandle_t handle;

    // Prepare input matrices
    float *A, *B, *C;
    int M, N, K;
    float alpha, beta;

    M = 3;    N = 4;    K = 7;
    alpha = 1.f;    beta = 0.f;

    // create cuBLAS handle
    cublasCreate(&handle);

    srand(2019);
    A = getMatrix(K, M);
    B = getMatrix(N, K);
    C = getMatrix(M, N);

    std::cout << "A:" << std::endl;
    printMatrix(A, K, M);
    std::cout << "B:" << std::endl;
    printMatrix(B, N, K);
    std::cout << "C:" << std::endl;
    printMatrix(C, M, N);

    // Gemm
    cublasSgemm(handle, CUBLAS_OP_T, CUBLAS_OP_T, 
        M, N, K, &alpha, A, K, B, N, &beta, C, M);

    cudaDeviceSynchronize();
    std::cout << "C out:" << std::endl;
    printMatrix(C, M, N);

    cublasDestroy(handle);
    cudaFree(A);    cudaFree(B);    cudaFree(C);
    return 0;

通过链接 cuBLAS 库使用nvcc编译代码：

$ nvcc -run -gencode arch=compute_70,code=sm_70 -lcublas -o cublasSgemm ./cublasSgemm.cpp

以下代码片段显示了执行的输出：

A:
 0.0492 0.4790 0.0226
 0.7750 0.2794 0.8169
 0.3732 0.6248 0.2636
 0.9241 0.5841 0.8532
 0.7188 0.5052 0.5398
 0.9869 0.6572 0.0520
 0.6815 0.7814 0.5988
B:
 0.8957 0.0481 0.7958 0.7825 0.3264 0.5189 0.5018
 0.4559 0.6342 0.0759 0.5051 0.1132 0.0985 0.2802
 0.3926 0.9153 0.6534 0.0174 0.1790 0.5775 0.6015
 0.0322 0.2963 0.1068 0.5720 0.2832 0.7640 0.6240
C:
 0.9647 0.5454 0.2229 0.8604
 0.5935 0.0186 0.6430 0.9198
 0.5375 0.1448 0.3757 0.1718
C out:
 1.1785 2.5682 2.4854 0.6066
 0.5817 0.8091 1.1724 2.0773
 2.0882 1.4503 2.1331 1.8450

如cublasSgemm()函数调用中所述，矩阵A和B是转置矩阵。我们将原始的主列大小传递给cublasSgemm()函数作为lda、ldb和ldc，我们可以看到操作按预期工作。

多 GPU 操作

cuBLAS 库的 cuBLAS-XT API 在多个 GPU 上运行时提供了 cuBLAS 的 3 级操作。使用此 API，您的应用程序可以使用多 GPU 计算操作。此片段显示了使用 cuBLAS-XT 的基本操作：

cublasXtHandle_t handle;
cublasXtCreate(&handle);

cudaGetDeviceCount(&num_of_total_devices);
devices = (int *)calloc(num_of_devices, sizeof(int));
for (int i = 0; i < num_of_devices; i++)
    devices[i] = i;
cublasXtDeviceSelect(handle, num_of_devices, devices);

cublasXtSgemm( ... );

cublasXtDestroy(handle);

cublasXtSgemm()接口与cublasSgemm()函数相同，因此我们可以轻松使用多个 GPU 的计算性能。例如，我们可以使用存储库中的示例代码在两个 GPU 上获得以下结果。这种性能可能会因您的 GPU 和系统配置而有所不同：

cuBLAS 库提供了许多多功能的线性代数操作。因此，您应该检查库中提供的必要功能。此外，您将需要一个如何使用该功能的示例。以下项目是文档和示例的链接。因此，建议您在需要基于 cuBLAS 实现应用程序时经常检查这两个文档：

英伟达的 cuBLAS 编程指南-参考指南：docs.nvidia.com/cuda/cublas/index.html
GPU 上的矩阵计算：CUBLAS、CUSOLVER 和 MAGMA 示例 by Andrzej Chrzȩszczyk and Jacob Anders: developer.nvidia.com/sites/default/files/akamai/cuda/files/Misc/mygpu.pdf

使用 cuBLAS 进行混合精度操作

cuBLAS 库支持混合精度计算。这种计算意味着使用不同精度进行操作，例如，使用单精度和半精度变量进行计算，或者使用单精度和字符(INT8)进行计算。当我们需要以降低精度获得更高性能，同时又获得更高准确性时，这种技术是有用的。

cuBLAS 库提供了cublasGemmEx()和cublas{S/C}gemmEx()来支持混合精度操作的 GEMM 运算。它们是cublas<t>gemm()的扩展，接受每个A、B和C矩阵的指定数据类型。以下表格显示了cublasGemmEx()的精度支持矩阵，以及 cuBLAS 库中其他可替换的 API：

计算类型	A 类型 / B 类型	C 类型	可替换的 API
`CUDA_R_16F`	`CUDA_R_16F`	`CUDA_R_16F`	`cublasHgemm()`
`CUDA_R_32I`	`CUDA_R_8I`	`CUDA_R_32I`	N/A
`CUDA_R_32F`	`CUDA_R_16F`	`CUDA_R_16F`	`cublasSgemmEx()`
`CUDA_R_8I`	`CUDA_R_32F`
`CUDA_R_16F`	`CUDA_R_32F`
`CUDA_R_32F`	`CUDA_R_32F`
`CUDA_R_64F`	`CUDA_R_64F`	`CUDA_R_64F`	`cublasDgemm()`
`CUDA_C_32F`	`CUDA_C_8I`	`CUDA_C_32F`	`cublasCgemmEx()`
`CUDA_C_32F`	`CUDA_C_32F`
`CUDA_C_64F`	`CUDA_C_64F`	`CUDA_C_64F`	`cublasZgemm()`

您可以看到cublasGemmEx()可以覆盖cublas{S/C}gemmEx()函数的操作。因此，我们将在本节中介绍cublasGemmEx()。

cublasGemmEx()函数的最后一个参数cublasGemmAlgo_t指定了矩阵乘法的算法。有了这个参数，我们可以选择是否使用 TensorCore。CUBLAS_GEMM_DEFAULT选择 GEMM 算法并在 CUDA 核心上运行。另一方面，CUBLAS_GEMM_DEFAULT_TENSOR_OP选择使用张量核心的算法。如果给定条件下 TensorCore 不可用，cuBLAS 会选择使用 CUDA 核心的算法。这种情况可能是没有张量核心的 GPU 或矩阵大小不符合张量核心操作的方式，即以四的倍数（** 4*）。

混合精度的 GEMM

现在，让我们尝试使用 cuBLAS GEMM 操作进行混合精度。实现后，我们将介绍矩阵大小如何影响操作。完全实现的版本在02_sgemm_mixed_precision/cublasGemmEx.cu中：

此代码使用自定义内存管理类CBuffer来简化混合精度和复制的处理，但也可以使用统一内存。对于 cuBLAS 操作，我们应该在代码中包含cublas_v2.h：

#include <cublas_v2.h>
#include "helper.cuh"    // for CBuffer and printMatrix()

现在，让我们实现main()函数。首先，我们将创建和初始化A、B和C矩阵。以下代码片段显示了如何使用CBuffer类并初始化矩阵：

int M = 4, N = 5, K = 6;
CBuffer<half> A, B;
CBuffer<float> C;

A.init(K * M, true);
B.init(N * K, true);
C.init(N * M, true);

指定A、B和C的精度类型，并测试各种精度，我们需要指定一些 CUDA 数据类型参数：

cudaDataType TYPE_A, TYPE_B, TYPE_C;
if (typeid(*A.h_ptr_) == typeid(float)) {
    TYPE_A = TYPE_B = CUDA_R_32F;
}
else if (typeid(*A.h_ptr_) == typeid(half)) {
    TYPE_A = TYPE_B = CUDA_R_16F;
}
else if (typeid(*A.h_ptr_) == typeid(int8_t)) {
    TYPE_A = TYPE_B = CUDA_R_8I;
}
else {
    printf("Not supported precision\n");
    return -1;
}

if (typeid(*C.h_ptr_) == typeid(float)) {
    TYPE_C = CUDA_R_32F;
}
else if (typeid(*C.h_ptr_) == typeid(int)) {
    TYPE_C = CUDA_R_32I;
}
else {
    printf("Not supported precision\n");
    return -1;
}

对于 cuBLAS 操作，我们应该初始化cublas_handle、alpha和beta：

float alpha = 1.f, beta = 0.f;
cublasHandle_t cublas_handle;
cublasCreate(&cublas_handle);

然后，我们将数据复制到 GPU：

A.cuda(true);
B.cuda(true);
C.cuda(true);

然后，我们调用cublasGemmEx()函数如下：

cublasGemmEx(cublas_handle,
                CUBLAS_OP_N, CUBLAS_OP_N,
                M, N, K,
                &alpha, A.d_ptr_, TYPE_A, M, B.d_ptr_, TYPE_B, K,
                &beta,  C.d_ptr_, TYPE_C, M, TYPE_C,
                CUBLAS_GEMM_DEFAULT);

要查看矩阵的值，我们可以使用helper.h中定义的printMatrix()：

std::cout << "A:" << std::endl;
printMatrix(A.h_ptr_, K, M);
std::cout << "B:" << std::endl;
printMatrix(B.h_ptr_, N, K);

printMatrix()使用函数重载方法定义，以允许以其他数据类型相同格式打印半精度值。部分定义如下：

template <typename T>
void printMatrix(const T *matrix, const int ldm, const int n) {
    std::cout << "[" << __FUNCTION__ << "]:: 
                Not supported type request" << std::endl;
}
void printMatrix(const float *matrix, const int ldm, const int n) {
    for (int j = 0; j < n; j++) {
        for (int i = 0; i < ldm; i++)
            std::cout << std::fixed << std::setw(8) << 
                         std::setprecision(4) << 
                         matrix[IDX2C(i, j, ldm)];
        std::cout << std::endl;
    }
}
void printMatrix(const half *matrix, const int ldm, const int n) {
    for (int j = 0; j < n; j++) {
        for (int i = 0; i < ldm; i++)
            std::cout << std::fixed << std::setw(8) <<
                         std::setprecision(4) << __half2float(matrix[IDX2C(i, j, ldm)]);
        std::cout << std::endl;
    }
}
... ( functions for other data types ) ...

然后，代码将对给定的A、B和C矩阵进行 GEMM 操作。以下是当M为4，N为5，M为6时的输出示例：

$ nvcc -run -m64 -std=c++11 -I/usr/local/cuda/samples/common/inc -gencode arch=compute_70,code=sm_70 -lcublas -o cublasGemmEx ./cublasGemmEx.cu
A:
 0.0049 0.0479 0.0023 0.0775 0.0279 0.0817
 0.0373 0.0625 0.0264 0.0924 0.0584 0.0853
 0.0719 0.0505 0.0540 0.0987 0.0657 0.0052
 0.0682 0.0781 0.0599 0.0896 0.0048 0.0796
B:
 0.0624 0.0965 0.0545 0.0223 0.0861
 0.0594 0.0019 0.0643 0.0920 0.0537
 0.0145 0.0376 0.0172 0.0221 0.0881
 0.0285 0.0319 0.0161 0.0677 0.0235
 0.0814 0.0695 0.0414 0.0392 0.0296
 0.0446 0.0688 0.0403 0.0018 0.0971
C:
 0.0509 0.0117 0.0877 0.0445 0.0830
 0.0742 0.0242 0.0136 0.0625 0.0681
 0.0362 0.0046 0.0265 0.0963 0.0638
 0.0070 0.0446 0.0516 0.0194 0.0089
C out:
 0.0153 0.0228 0.0143 0.0292 0.0113
 0.0200 0.0118 0.0214 0.0081 0.0138
 0.0098 0.0168 0.0132 0.0199 0.0125
 0.0269 0.0120 0.0222 0.0085 0.0228

现在，让我们尝试其他数据类型，并查看cublasGemmEx()如何对给定的矩阵进行操作。提供的示例还输出了操作的执行时间以衡量性能：

如果矩阵A或矩阵B是转置矩阵，我们应该修改什么？
是否有任何更适合的矩阵大小进行操作？通过改变大小来比较执行时间。
每种数据类型是否有任何更适合的矩阵大小？如果尝试INT8精度，会出现错误。如何解决这个问题？改变大小并查看cublasGemmEx()如何支持INT8操作。

TensorCore 的 GEMM

TensorCore 提供了张量点运算的加速性能。它支持 Volta 架构中的 FP16，Turing 架构中的INT8和INT4。因此，我们应该使用降低的精度或混合精度来使用 TensorCore。

之前我们使用CUBLAS_GEMM_DEFAULT作为 cuBLAS GEMM 算法，它在操作中使用 CUDA 核心。要使用 TensorCore，我们应该使用CUBLAS_GEMM_DEFAULT_TENSOR_OP。为了利用 TensorCore，您的操作数矩阵的每个维度都应该是 4 的倍数。这是 TensorCore 的WMMA（Warp Matrix Multiply Accumulate的缩写）操作优化的单位大小。例如，矩阵A（8,192×8,192）和B（8,192×8,192）的矩阵乘法与A（8,192×8,192）和B（8,192×8,190）的操作相比，性能要高得多。您也可以通过性能分析来确认这个操作。

以下时间线是使用矩阵A（8,192×8,192）和矩阵B（8,192×8,190）进行矩阵乘法的结果：

此外，这个时间线图像是矩阵A(8,192×8,192)和矩阵B(8,192×8,192)的矩阵乘法的结果：

两个测试都在 CUDA C/C++中使用CUBLAS_GEMM_DEFAULT_TENSOR_OP，但是使用 TensorCore 的 GEMM 操作比使用 CUDA 核心快 6.7 倍。由于 TensorCore 基于矩阵大小可用，nvcc将使用特殊的内核函数编译代码，从volta_s884g开始。总之，如果要获得 TensorCore 的好处，可以将矩阵填充到 4 的倍数。这可能会增加开销，但是 TensorCore 的性能收益可能会超过开销。

NVIDIA 在其开发博客网站提供了如何使用 cuBLAS 库编程 TensorCores 的方法(devblogs.nvidia.com/programming-tensor-cores-cuda-9)。本文档还介绍了其他可用的方法。但是，使用 cuBLAS 库可以为您提供最快的性能，正如来自奥克岭国家实验室的一篇论文所证明的那样——NVIDIA Tensor Core 可编程性、性能和精度 (arxiv.org/pdf/1803.04014.pdf)。

用于并行随机数生成的 cuRAND

许多应用程序在模拟或概率分析中使用伪随机数。尽管它的常规用途，但大量的随机数生成过程需要很长时间。一个解决方案是并行生成随机数，但每个多线程应该有不同的随机种子，以便独立生成随机数。

cuRAND 库使 GPU 能够从 GPU 生成大量随机数。此库可从主机代码或设备代码中使用。主机 API 仅允许使用主机代码生成随机数。因此，您可以直接将生成的数据用于其他内核函数。设备 API 允许在内核代码中生成随机数，因此您可以在执行过程中创建具有自己随机生成数的 CUDA 线程。

cuRAND 主机 API

首先，您需要使用curandGenerator()创建所需类型的新生成器。然后，设置所需种子和顺序的生成器选项。例如，您可以使用curandSetPseudoRandomGeneratorSeed()生成伪随机生成器：

curandGenerator_t curand_gen;
curandCreateGenerator(&curand_gen, CURAND_RNG_PSEUDO_DEFAULT);
curandSetPseudoRandomGeneratorSeed(curand_gen, 2019UL);

然后，您可以使用curandGenerate()生成随机数。有九种不同的生成函数。例如，您可以使用curandGenerateUnifrom()生成均匀分布的浮点值：

curandGenerateUniform(curand_gen, random_numbers_on_device_memory, length);

cuRAND 编程指南提供了各种生成函数的描述：docs.nvidia.com/cuda/curand/host-api-overview.html#generation-functions。

在使用随机数生成后，您可以使用销毁函数终止 cuRAND 生成器：

curandDestroyGenerator(curand_gen);

现在，让我们实现一个使用几个 cuRAND API 生成随机数的应用程序。完全实现的版本是03_curand/curand_host.cpp。因此，您可以根据需要修改代码并测试其他函数。

首先，我们应该包括curand.h用于 cuRAND 主机 API 和其他与 CPP 相关的头文件，如下所示：

#include <iostream>
#include <iomanip>
#include <curand.h>

假设我们将创建一个用随机数初始化的矩阵。我们需要实现printMatrix()函数，以便查看生成的随机数，如下所示：

#define IDX2C(i, j, ld) (((j) * (ld)) + (i))

template <typename T>
void printMatrix(const T *matrix, const int ldm, const int n)
{
    for (int j = 0; j < ldm; j++) {
        for (int i = 0; i < n; i++)
            std::cout << std::fixed << std::setw(12) 
                    << std::setprecision(4) << matrix[IDX2C(i, j, ldm)];
        std::cout << std::endl;
    }
}

然后，我们将按以下方式分配所需的内存空间。现在，我们将实现main()函数，该函数初始化随机数并使用printMatrix()打印结果。首先，我们将为操作初始化 cuRAND 句柄，如下所示：

// create curand generator & set random seed
curandGenerator_t curand_gen;
curandCreateGenerator(&curand_gen, CURAND_RNG_PSEUDO_DEFAULT);
curandSetPseudoRandomGeneratorSeed(curand_gen, 2019UL);

您可以根据需要更改随机种子。接下来要做的是分配内存空间。为了简化操作的评估，我们将使用统一内存，因为 cuRAND 函数将在 GPU 上生成随机数：

size_t size = M * N;
unsigned int *np_random;
float *fp_random;
cudaMallocManaged((void**)&np_random, sizeof(*np_random) * size);
cudaMallocManaged((void**)&fp_random, sizeof(*fp_random) * size);

接下来，我们将在给定的内存空间中生成随机数。我们将使用整数内存空间（np_random）进行随机数生成，使用浮点内存空间（fp_random）进行均匀分布的随机数生成，如下所示：

// random number generation
std::cout << "Generated random numbers" << std::endl;
curandGenerate(curand_gen, np_random, size);
cudaDeviceSynchronize();
printMatrix(np_random, M, N);

// uniform distributed random number generation
std::cout << "Generated uniform random numbers" << std::endl;
curandGenerateUniform(curand_gen, fp_random, size);
cudaDeviceSynchronize();
printMatrix(fp_random, M, N);

因为我们使用统一内存，我们可以允许 GPU 和主机共享相同的内存地址，并且可以通过同步它们来查看输出值。最后，我们可以终止 cuRAND 句柄和内存，如下所示：

// terminates used resources
curandDestroyGenerator(curand_gen);
cudaFree(np_random);
cudaFree(fp_random);

现在，是时候编译和运行代码了。使用 cuRAND API 编译代码应该为nvcc编译器提供-lcurand。当M = 3和N = 5时，输出如下：

$ nvcc -run -gencode arch=compute_70,code=sm_70 -lcurand -o curand_host curand_host.cpp
Generated random numbers
 3395652512 793372546 2133571103 595847267 2461872808
 595847267 2461872808 500895635 498154070 2385617847
 498154070 2385617847 196336856 388563169 745758309
Generated uniform random numbers
 0.7134 0.0830 0.1458 0.2099 0.6066
 0.2099 0.6066 0.3078 0.5122 0.8856
 0.5122 0.8856 0.3530 0.8477 0.8370

我们已经介绍了 CUDA 如何使用主机 API 生成随机数，但在某些情况下，最好设计 CUDA 内核来生成随机数。我们称之为设备 API，我们可以从每个 CUDA 线程获取随机数。

cuRAND 设备 API

使用设备 API，我们可以在 CUDA 设备上设置生成器种子并生成随机数。

首先，我们需要准备一个curandState_t的设备内存空间，以便并行地为 CUDA 线程提供随机种子。这可以像正常的设备内存分配代码一样完成，如下所示：

cudaMalloc((void **)&devStates, length * sizeof(curandState_t));

在您的内核代码中，我们需要使用curand_init()初始化随机种子。这个函数需要种子、序列号和偏移量。然后，这个函数设置状态。对于相同的种子，cuFFT 总是生成相同的状态。要生成随机值，使用curand()函数。与主机的生成函数一样，设备 API 有各种生成函数。例如，均匀分布的随机数生成可以这样做：

int idx = blockIdx.x * blockDim.x + threadIdx.x;
curand_init(2019UL, idx, 0, &state[idx]);
generated_out[idx] = curand_uniform(&state[idx]);

cuRAND 库提供了各种数据类型和随机分布的生成函数。要找到您需要的生成函数，请查看 cuRAND 开发人员指南的设备 API 概述。在随机数生成之后，设备状态缓冲区应该像正常内存一样终止，如下所示：

cudaFree(devStates);

现在，我们将创建一个使用 cuRAND 设备 API 的应用程序。完全实现的代码是curand_device.cu，所以您也可以修改和测试代码。首先，我们应该包括curand_kernel.h文件和其他 C++所需的头文件，如下所示：

#include <iostream>
#include <iomanip>
#include <curand_kernel.h>

我们将编写setup_kernel()，为每个 CUDA 线程初始化一个随机种子，如下所示：

  __global__ void setup_kernel(curandState_t *state)
  {
      int idx = blockIdx.x * blockDim.x + threadIdx.x;
      // Each thread gets same seed, 
      // a different sequence number, no offset */
      curand_init(2019UL, idx, 0, &state[idx]);
  }

编写两个随机数生成函数：generate_kernel()和generate_uniform_kernel()。我们将生成一个 32 位整数和一个均匀分布的单精度浮点数随机数：

__global__ void generate_kernel(unsigned int *generated_out, 
                                curandState_t *state)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    generated_out[idx] = curand(&state[idx]) & 0xFF;
}

__global__ void generate_uniform_kernel(float *generated_out, 
                                        curandState_t *state)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    generated_out[idx] = curand_uniform(&state[idx]);
}

现在，我们将实现main()函数并初始化设备状态缓冲区：

cudaMalloc((void **)&devStates, sizeof(curandState) * size);
setup_kernel<<<(size + BLOCK_DIM - 1) / BLOCK_DIM, BLOCK_DIM>>>(devStates);

然后，使用generate_kernel()生成随机数。为了方便起见，我们将使用统一内存空间，并验证来自主机的输出。之后，我们将打印结果，如下所示：

// random number generation
std::cout << "Generated random numbers" << std::endl;
cudaMallocManaged((void**)&np_random, sizeof(*np_random) * size);
generate_kernel<<<(size + BLOCK_DIM - 1) / BLOCK_DIM, BLOCK_DIM>>>
                (np_random, const_cast<curandState_t *>(devStates));
cudaDeviceSynchronize();
printMatrix(np_random, M, N);

同样地，我们将使用generate_uniform_kernel()创建均匀分布的随机数，如下所示：

// uniform distributed random number generation
std::cout << "Generated uniform random numbers" << std::endl;
cudaMallocManaged((void**)&fp_random, sizeof(*fp_random) * size);
generate_uniform_kernel<<<(size + BLOCK_DIM - 1) / BLOCK_DIM, BLOCK_DIM>>>
                (fp_random, const_cast<curandState_t *>(devStates));
cudaDeviceSynchronize();
printMatrix(fp_random, M, N);

// terminates used resources
curandDestroyGenerator(curand_gen);
cudaFree(np_random);
cudaFree(fp_random);

现在，是时候编译和运行代码了。为了使用 cuRAND 编译代码，API 应该为nvcc编译器提供-lcurand。当M等于3，N等于5时，输出如下：

$ nvcc -run -gencode arch=compute_70,code=sm_70 -lcurand -o curand_device curand_device.cpp
Generated random numbers
 3395652512 793372546 2133571103 595847267 2461872808
 595847267 2461872808 500895635 498154070 2385617847
 498154070 2385617847 196336856 388563169 745758309
Generated uniform random numbers
 0.8064 0.2783 0.2971 0.2386 0.7491
 0.2386 0.7491 0.4782 0.1060 0.2922
 0.1060 0.2922 0.1823 0.6199 0.9137

当您比较主机 API 和设备 API 的输出数字时，生成的随机数是相同的，而均匀随机数不是。如果在第二次随机数生成之前重置随机种子，这个问题可以得到解决。

cuRAND 与混合精度 cuBLAS GEMM

以前，我们使用 C++随机数生成器来初始化 GEMM 操作的矩阵。当我们想要生成随机数时，这个函数很方便。然而，您可能会发现，在上一节中生成大量随机数需要很长时间。在本节中，我们将介绍 cuRAND API 如何与 cuBLAS GEMM 操作配合使用。完全实现的版本是gemm_with_curand_host.cpp文件。让我们看看这是如何实现的：

目前，cuRAND 库中没有低精度的随机数生成器。此外，我们需要将半精度数转换为浮点数以评估输出。因此，我们需要在 GPU 上创建类型转换函数，如下所示：

namespace fp16{
__global__ void float2half_kernel(half *out, float *in)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    out[idx] = __float2half(in[idx]);
}

void float2half(half *out, float *in, size_t length)
{
    float2half_kernel<<< (length + BLOCK_DIM - 1) / BLOCK_DIM, BLOCK_DIM >>>(out, in);
}

现在，我们将编写一个使用 cuRAND 主机 API 的随机数生成函数。正如我们之前讨论的，当我们需要使用半精度数据时，我们应该将生成的随机数从浮点数转换为半精度。这个函数可以实现如下：

template <typename T>
typename std::enable_if<(std::is_same<T, float>::value), float>::type
*curand(curandGenerator_t generator, size_t length)
{
    T *buffer = nullptr;
    cudaMalloc((void **)&buffer, length * sizeof(float));
    curandGenerateUniform(generator, buffer, length);
    return buffer;
}
template <typename T>
typename std::enable_if<std::is_same<T, half>::value, half>::type
*curand(curandGenerator_t generator, size_t length)
{
    T *buffer = nullptr;
    float *buffer_fp32;

    cudaMalloc((void **)&buffer_fp32, length * sizeof(float));
    curandGenerateUniform(generator, buffer_fp32, length);

    cudaMalloc((void **)&buffer, length * sizeof(T));
    fp16::float2half(buffer, buffer_fp32, length);
    cudaFree(buffer_fp32);

    return buffer;
}

在main()函数中定义一些控制 GEMM 操作的本地变量：

void *d_A, *d_B, *d_C;
cudaDataType AType, BType, CType, computeType;
int M = 8192, N = 8192, K = 8192;
float alpha = 1.f, beta = 1.f;
std::string precision = "fp32";
bool tensor_core = true;

在这段代码中，我们确定了 GEMM 操作的大小、数据类型和操作类型。

现在，让我们创建输入缓冲区数组，并设置参数，以及操作精度：

if (precision == "fp32") {
    auto *a = curand<float>(curand_gen, M * K);
    auto *b = curand<float>(curand_gen, K * N);
    auto *c = curand<float>(curand_gen, M * N);
    AType = BType = CType = CUDA_R_32F;
    computeType = CUDA_R_32F;
    d_A = a, d_B = b, d_C = c;
}
else if (precision == "fp16") {
    auto *a = curand<half>(curand_gen, M * K);
    auto *b = curand<half>(curand_gen, K * N);
    auto *c = curand<float>(curand_gen, M * N);
    AType = BType = CUDA_R_16F, CType = CUDA_R_32F;
    computeType = CUDA_R_32F;
    d_A = a, d_B = b, d_C = c;
}
else {
    exit(EXIT_FAILURE);
}

创建 cuRAND 和 cuBLAS 句柄如下：

cublasCreate(&cublas_handle);
curandCreateGenerator(&curand_gen, CURAND_RNG_PSEUDO_DEFAULT);
curandSetPseudoRandomGeneratorSeed(curand_gen, 2019UL);

然后，我们应该确定操作类型以使用 TensorCores：

cublasGemmAlgo_t gemm_algo = (tensor_core) ? 
                             CUBLAS_GEMM_DEFAULT_TENSOR_OP : CUBLAS_GEMM_DEFAULT;

然后，我们可以调用cublasGemmEx()函数，提供 FP32 和 FP16 操作，如下所示：

cublasGemmEx(cublas_handle, CUBLAS_OP_N, CUBLAS_OP_N,
             M, N, K,
             &alpha, d_A, AType, M, d_B, BType, K,
             &beta,  d_C, CType, M,
             computeType, gemm_algo);

与以前的版本相比，GEMM 操作应该表现出类似的性能。但是，您可能会发现整个应用程序的速度得到了提高，因为 GPU 上的并行随机数生成比主机上的生成要快得多。

cuRAND 开发人员指南将帮助您找到其他随机数生成器、选项和分布。该文档位于docs.nvidia.com/pdf/CURAND_Library.pdf。

用于 GPU 中的快速傅里叶变换的 cuFFT

cuFFT 库为FFT（快速傅里叶变换）算法提供了 GPU 加速操作。程序员可以利用 GPU 计算能力转换实数或复数数据，并对转换后的信号应用 GPU 核操作。此外，支持的函数与 FFTW 库匹配，因此我们可以将主机项目迁移到 GPU。

为了处理 FFT 样本的维度信息，需要使用 cuFFT 来使用cufftPlan1D()、cufftPlan2D()或cufftPlan3D()来创建一个计划句柄。如果样本数据具有批处理和步幅布局，我们应该使用cufftPlanMany()。如果样本大小大于 4GB，我们应该使用64作为计划函数的后缀来支持该大小。例如，cufftPlanMany64()支持cufftPlanMany()函数之上的更大样本。

cuFFT 库支持多 GPU 操作。首先，您需要使用cufftCreate()创建一个空计划。然后，我们可以使用cufftXtSetGPUs()指定将执行操作的 GPU 列表。之后，我们可以使用先前介绍过的普通计划生成函数生成一个计划。以下表格显示了计划生成函数的类别：

基本计划	多 GPU 计划
简单计划	`cufftPlan{1d,2d,3d}()`
高级数据布局	`cufftPlanMany()`
FP16 操作	`cufftXtMakePlanMany()`

然后，您可以使用cufftExec()函数对样本数据进行前向（FFT）和反向（IFFT）变换。cuFFT 库提供三种数据转换：复杂到复杂、实到复杂和复杂到实。其操作数据类型可以是浮点数或双精度数。

变换方向	浮点数	双精度
复杂到复杂	`cufftExecC2C()` `cufftXtExecDescriptorC2C()`	`cufftExecZ2Z()` `cufftXtExecDescriptorZ2Z()`
实到复	`cufftDExecR2C()` `cufftXtExecDescriptorR2C()`	`cufftExecD2Z()` `cufftXtExecDescriptorD2Z()`
复到实	`cufftExecC2R()` `cufftXtExecDescriptorC2R()`	`cufftExecZ2D()` `cufftXtExecDesciptorZ2D()`
所有	`cufftXtExec()` / `cufftXtExecDesciptor()`

cuFFT 操作要么是正向要么是反向，操作应该与另一个方向配对。

在函数名称中，将实数据和复杂数据之间进行转换的函数，例如R2C和C2R，具有隐含的方向信息。这个特性可以帮助您避免必须进行额外的操作来将您的数据从实域转换为复杂数据类型。同时，由于每个计划都有变换方向信息，因此您必须创建一个额外的计划。

另一方面，您必须为复到复变换提供变换方向信息，例如C2C和Z2Z。对于反演操作，您不必创建另一个 cuFFT 句柄，因为计划应该是相同数据类型的操作。

cufftXtExec()和cufftXtExecDescriptor()函数可以对任何给定的数据类型执行变换，因为在创建 cuFFT 计划时应该提供每个输入数据的数据类型信息。

cuFFT 的基本用法

现在让我们尝试使用 cuFFT。完全实现的版本是04_cufft/cufft.1d.cpp文件。让我们讨论一下它是如何实现的：

首先，从一些头文件开始：C++，CUDA，cuRAND 和 cuFFT：

#include <iostream>
#include <iomanip>
#include <cuda_runtime.h>
#include <cufft.h>
#include <curand.h>

在这个 FFT 操作中，我们将进行实到复和复到实的转换。因此，让我们声明一些自定义数据类型，Real和Complex，以简化代码。可以这样做：

typedef cufftReal Real;
typedef cufftComplex Complex;

现在，让我们从main()函数开始。对于输入的样本数据，我们将使用统一内存以便简化主机和 GPU 之间的数据传输。转换后的数据只能在 GPU 上使用。因此，可以分配内存空间如下：

cudaMallocManaged((void**)&p_sample, sizeof(Real) * sample_size * batch_size);
cudaMalloc((void**)&d_freq, sizeof(Complex) * sample_size * batch_size);

然后，我们将使用 cuRAND 主机 API 来初始化输入数据如下：

curandGenerator_t curand_gen;
curandCreateGenerator(&curand_gen, CURAND_RNG_PSEUDO_DEFAULT);
curandSetPseudoRandomGeneratorSeed(curand_gen, 2019UL);
curandGenerateUniform(curand_gen, p_sample, sample_size * batch_size);

而且，我们应该初始化 cuFFT 计划以进行正向和反向变换。由于它们具有不同的数据类型，我们应该分别为实到复和复到实的转换创建两个计划：

int rank = 1;
int stride_sample = 1, stride_freq = 1;
int dist_sample = sample_size, dist_freq = sample_size / 2 + 1;
int embed_sample[] = {0};
int embed_freq[] = {0};
cufftPlanMany(&plan_forward, rank, &sample_size,
                             embed_sample, stride_sample, 
                             dist_sample, 
                             embed_freq, stride_freq, dist_freq,
                             CUFFT_R2C, batch_size);
cufftPlanMany(&plan_inverse, rank, &sample_size,
                             embed_freq, stride_freq, dist_freq, 
                             embed_sample, stride_sample, 
                             dist_sample,
                             CUFFT_C2R, batch_size);

现在，我们可以使用给定的 cuFFT 计划进行正向或反向变换。为了测量执行时间，我们可以使用 CUDA 事件来包含这些操作：

cufftExecR2C(plan_forward, p_sample, d_freq);
cufftExecC2R(plan_inverse, d_freq, p_sample);

然后，我们可以使用以下命令编译代码：

$ nvcc -run -gencode arch=compute_70,code=sm_70 -lcufft -lcurand -o cufft.1d cufft.1d.cpp

cufft.1d命令将报告每个步骤的变换时间如下：

FFT operation time for 1048576 elements with 512 batch..
Forward (ms): 21.5322
Inverse (ms): 21.4

混合精度的 cuFFT

cuFFT 库提供了扩展的 CUDA 计算功能，例如 FP16 FFT 操作。完整版本是cufft.half.cpp文件。让我们讨论它的实现。

在这段代码中，我们应该使用cufftXtMakePlanMany()来创建计划，使用cufftXtExec()函数进行变换。cufftXtMakePlanMany()允许传递输入和输出数据类型，如果它们是 FP16 或 FP32。此外，我们应该为正向和反向变换创建两个计划，以涵盖实到复和复到实的转换。对于空的 cuFFT 计划，cufftXtMakePlanMany()可以指定样本大小、输入数据格式和类型、批处理大小等。例如，计划的创建可以实现如下：

int rank = 1;
int stride_sample = 1, stride_freq = 1;
long long int dist_sample = sample_size, dist_freq = sample_size / 2 + 1;
long long embed_sample[] = {0};
long long embed_freq[] = {0};
size_t workSize = 0;
cufftCreate(&plan_forward);
cufftXtMakePlanMany(plan_forward, 
        rank, &sample_size, 
        embed_sample, stride_sample, dist_sample, CUDA_R_16F, 
        embed_freq, stride_freq, dist_freq, CUDA_C_16F, 
        batch_size, &workSize, CUDA_C_16F);
cufftCreate(&plan_inverse);
cufftXtMakePlanMany(plan_inverse,
        rank, &sample_size,
        embed_freq, stride_freq, dist_freq, CUDA_C_16F,
        embed_sample, stride_sample, dist_sample, CUDA_R_16F,
        batch_size, &workSize, CUDA_R_16F);

在这个实现中，我们还必须考虑是否提供半精度的输入数据。您可以使用主机随机函数并将它们转换为半精度数据，但是，这段代码向您展示了如何使用 cuRAND 主机 API 来实现此目的，如下所示：

template <typename T>
typename std::enable_if<std::is_same<T, half>::value>::type
curand(curandGenerator_t generator, T *buffer, size_t length) {
    float *buffer_fp32;

    cudaMalloc((void **)&buffer_fp32, length * sizeof(float));
    curandGenerateUniform(generator, buffer_fp32, length);

    // convert generated single floating to half floating
    fp16::float2half(buffer, buffer_fp32, length);
    cudaFree(buffer_fp32);
}

因此，我们可以为 FFT 提供半精度的均匀分布随机数，并且我们可以使用cufftXtExec()进行正向和反向变换。变换性能如下：

 FFT operation time for 1048576 elements with 512 batch..
Forward (ms): 15.3236
Inverse (ms): 15.4881

多 GPU 的 cuFFT

cuFFT 的另一个用途是使用多个 GPU 进行大型 FFT 操作。为了做到这一点，我们必须使用cufftCreate()创建一个空的 cuFFT 计划，并使用cufftXtSetGPUs()提供 GPU 数量。例如，可以这样做：

cufftHandle cufft_plan;
int n_gpu = 2, devices[2] = {0,1};
cufftCreaet(&cufft_plan); // create an empty plan
cufftXtSetGPUs(cufft_plan, n_gpu, devices); // set multi-gpu information

GPU 的总数可能会根据系统而有所不同。现在，我们可以使用cufftXtMakePlanMany()生成 cuFFT 计划，以指定样本信息。例如，可以这样调用cufftXtMakePlanMany()：

size_t *work_size = (size_t*) new size_t[num_gpus];
cufftXtMakePlanMany(cufft_plan, 1 &sample_size, 
                    nullptr, 1, 1, CUDA_C_32F, 
                    nullptr, 1, 1,  CUDA_C_32F, 
                    batch_size, work_size, CUDA_C_32F);

cuFFT 库提供了cufftXtMalloc()，它为目标 GPU 准备 GPU 内存空间。然后，我们可以使用cufftXtMemcpy()函数将我们的数据复制到分配的内存中。例如，可以这样实现：

cudaLibXtDesc *d_sample;
cufftXtMalloc(cufft_plan, &d_sample, CUFFT_XT_FORMAT_INPLACE);
cufftXtMemcpy(cufft_plan, d_sample, h_sample, CUFFT_COPY_HOST_TO_DEVICE);

然后，我们可以使用cufftXtExecDesciptor()函数在多个 GPU 上执行 FFT。

cufftXtExecDesciptor(cufft_plan, d_sample, d_sample, CUFFT_FORWARD);

使用nvidia-smi，我们可以监视分布式内存分配和跨 GPU 的执行。经过的时间可能会有所不同，这取决于你的 GPU 和系统配置。

如果你想了解更多关于 cuFFT 库及其函数的信息，cuFFT 库用户指南(docs.nvidia.com/cuda/cufft/index.html)是一个很好的参考资料。

CUDA 示例代码是学习如何使用 cuFFT 函数的另一个很好的参考。示例代码放在NVIDIA_CUDA-10.x_Samples/7_CUDALibraries/CUFFT*目录中。你可以学习如何使用 CUDA 核心代码应用滤波操作，以及通过 cuFFT 的正向/反向变换来实现。

NPP 用于 GPU 图像和信号处理

NPP（NVIDIA 性能基元）库是一个默认的 CUDA 库，其中包含一组 GPU 加速处理函数，专注于图像和视频处理。虽然它能够在这些领域灵活开发，但开发人员可以节省应用程序开发时间。

NPP 库有两个功能部分：图像处理 API 和信号处理 API。图像处理 API 包括与图像滤波、压缩/解压缩、颜色转换、调整大小、颜色转换、统计操作等相关的工具。信号处理 API 包括滤波、转换等。你可以访问 NPP 的文档(docs.nvidia.com/cuda/npp)，查看它的配置和完整的功能列表。

CUDA 提供了许多基于 NPP 的示例。在本节中，我们将介绍 NPP 库的基本用法并讨论其应用。

使用 NPP 进行图像处理

首先，我们将介绍 NPP 库如何简化图像处理任务。在这之前，我们应该安装 FreeImage 库，以便能够轻松地加载和写入 JPEG 压缩图像文件。有三种选项可以用来准备库：

从 Ubuntu 存档中安装：

$ sudo apt-get install libfreeimage-dev

从源代码构建和安装：

$ wget http://downloads.sourceforge.net/freeimage/FreeImage3180.zip
$ unzip FreeImage3180.zip
$ cd FreeImage && make -j && sudo make install

使用已经安装了 CUDA Toolkit 的库。CUDA 示例代码中的 NPP 示例代码7_CUDALibraries/freeImageInteropNPP使用了 FreeImage 库。对于这个示例，NPP 头文件和库文件安装在 CUDA 示例目录中的7_CUDALibrires/common/FreeImage。如果你不想在你的机器上安装其他二进制文件，你可以使用这个。

现在，让我们实现基于 NPP 的图像处理应用程序。完全实现的代码是05_npp/imageFilter.cpp。这个文件以头文件开始：

#include <iostream>
#include <iomanip>
#include <cassert>
#include <cstring>
#include <cuda_runtime.h>
#include <npp.h>
#include <FreeImage.h>
#include <helper_timer.h>

在这个应用程序中，它有ImageInfo_t结构来方便地管理图像信息和数据：

struct ImageInfo_t
{
    /* image information */
    FIBITMAP* dib; // FreeImage bitmap
    int nHeight;   // image height size
    int nWidth;    // image width size
    int nPitch;    // image pitch size
    int nBPP;      // Bit Per Pixel (i.e. 24 for BGR color)
    int nChannel;  // number of channels 
    BYTE* pData;   // bytes from freeimage library

    /* CUDA */
    Npp8u *pDataCUDA; // CUDA global memory for nppi processing
    int nPitchCUDA;   // image pitch size on CUDA device
};

编写LoadImage()函数以加载 JPEG 图像。FreeImage库支持任何其他图像格式，所以你可以根据需要尝试其他图像。然后，我们将用加载的图像数据填充源图像信息管理结构。loadImage()函数的实现如下：

void LoadImage(const char *szInputFile, ImageInfo_t &srcImage) {
    FIBITMAP *pSrcImageBitmap = FreeImage_Load(FIF_JPEG, szInputFile, JPEG_DEFAULT);
    if (!pSrcImageBitmap) {
        std::cout << "Couldn't load " << szInputFile << std::endl;
        FreeImage_DeInitialise();
        exit(1);
    }

    srcImage.dib = pSrcImageBitmap;
    srcImage.nWidth = FreeImage_GetWidth(pSrcImageBitmap);
    srcImage.nHeight = FreeImage_GetHeight(pSrcImageBitmap);
    srcImage.nPitch = FreeImage_GetPitch(pSrcImageBitmap);
    srcImage.nBPP = FreeImage_GetBPP(pSrcImageBitmap);
    srcImage.pData = FreeImage_GetBits(pSrcImageBitmap);
    assert(srcImage.nBPP == (unsigned int)24); // BGR color image
    srcImage.nChannel = 3;
}

然后，编写一些 NPPI 辅助函数，从图像结构中提供 NPPI 图像大小和 NPPI ROI 大小数据，如下所示：

NppiSize GetImageSize(ImageInfo_t imageInfo)
{
    NppiSize imageSize;

    imageSize.width = imageInfo.nWidth;
    imageSize.height = imageInfo.nHeight;

    return imageSize;
}

NppiRect GetROI(ImageInfo_t imageInfo)
{
    NppiRect imageROI;

    imageROI.x = 0;    imageROI.y = 0;
    imageROI.width = imageInfo.nWidth;
    imageROI.height = imageInfo.nHeight;

    return imageROI;
}

然后，让我们实现基于 NPPI 的图像调整大小函数。在此函数中，我们将使用一开始讨论的nppiResize_8u_C3R()。NPP API 有命名约定规则，以明确说明它们的操作。根据其功能类别，它们的命名以nppi开头用于图像处理，以npps开头用于信号处理。例如，一个 NPP 图像处理函数nppiResize_8u_C3R()以nppi前缀开头，它将具有三个通道的无符号字符数据类型的输入数据调整大小到给定的 ROI（您可以在文档中了解更多关于这种约定的细节）：

int ResizeGPU(ImageInfo_t &dstImage, ImageInfo_t &srcImage, 
                 NppiSize &dstSize, NppiRect &dstROI, 
                 NppiSize &srcSize, NppiRect &srcROI, scale)
{
    // update output image size
    dstSize.width = dstROI.width = dstImage.nWidth;
    dstSize.height = dstROI.height = dstImage.nHeight;

    nppiResize_8u_C3R(srcImage.pDataCUDA, srcImage.nPitchCUDA, 
                      srcSize, srcROI, 
                      dstImage.pDataCUDA, dstImage.nPitchCUDA, 
                      dstSize, dstROI,
                      NPPI_INTER_LANCZOS);
    return 0;
}

为了与 CPU 的性能进行比较，我们将使用 FreeImage 的一个函数，如下所示：

void ResizeCPU(const char* szInputFile, ImageInfo_t &dstImage) {
    FreeImage_Rescale(dib, dstImage.nWidth, dstImage.nHeight, FILTER_LANCZOS3);
}

现在，让我们实现main()函数。首先，我们应该初始化 FreeImage 库并加载一个图像：

FreeImage_Initialise();
ImageInfo_t srcImage, dstImage;
LoadImage(szInputFile, srcImage);

然后，我们将初始化输入图像的 GPU 内存空间，如下所示。在此过程中，我们将使用 NPPI 函数初始化全局内存空间，并使用cudaMemcpy2D()将加载的图像传输到全局内存中：

// copy loaded image to the device memory
srcImage.pDataCUDA = 
             nppiMalloc_8u_C3(srcImage.nWidth, srcImage.nHeight, 
                              &srcImage.nPitchCUDA);
cudaMemcpy2D(srcImage.pDataCUDA, srcImage.nPitchCUDA, 
             srcImage.pData, srcImage.nPitch, 
             srcImage.nWidth * srcImage.nChannel * sizeof(Npp8u), 
             srcImage.nHeight,
             cudaMemcpyHostToDevice);

之后，我们将初始化输出内存空间，并提供调整后的图像大小信息，如下所示：

std::memcpy(&dstImage, &srcImage, sizeof(ImageInfo_t));
dstImage.nWidth *= scaleRatio;
srcImage.nHeight *= scaleRatio;
dstImage.pDataCUDA = 
                nppiMalloc_8u_C3(dstImage.nWidth, dstImage.nHeight, 
                                 &dstImage.nPitchCUDA);

然后，我们调用已经实现的ResizeGPU()和ResizeCPU()函数。对于每个操作，我们将使用cudaEvent来测量 GPU 上的执行时间：

RunNppResize(dstImage, srcImage, dstImageSize, dstROI, srcImageSize, srcROI, scaleRatio);
RunCpuResize(szInputFile, dstImage);

为了验证，我们将结果保存到文件中。为此，我们应该创建一个 FreeImage 位图，并将调整大小后的图像复制到内存空间中。然后，我们可以保存输出图像，如下所示：

// Save resized image as file from the device
FIBITMAP *pDstImageBitmap = 
                FreeImage_Allocate(dstImage.nWidth, dstImage.nHeight, 
                                   dstImage.nBPP);

dstImage.nPitch = FreeImage_GetPitch(pDstImageBitmap);
dstImage.pData = FreeImage_GetBits(pDstImageBitmap);

cudaMemcpy2D(dstImage.pData, dstImage.nPitch, 
             dstImage.pDataCUDA, dstImage.nPitchCUDA, 
             dstImage.nWidth * dstImage.nChannel * sizeof(Npp8u),
             dstImage.nHeight, cudaMemcpyDeviceToHost);

FreeImage_Save(FIF_JPEG, pDstImageBitmap, szOutputFile, JPEG_DEFAULT);

之后，我们最终可以终止相关资源：

nppiFree(srcImage.pDataCUDA);
nppiFree(dstImage.pDataCUDA);

FreeImage_DeInitialise();

使用链接的 NPP 和 FreeImage 库使用nvcc编译代码：

$ nvcc -run -m64 -std=c++11 -I/usr/local/cuda/samples/common/inc -gencode arch=compute_70,code=sm_70 -lnppc -lnppif -lnppisu -lnppig -lnppicom -lnpps -lfreeimage -o imageFilter ./imageFilter.cpp

因此，当比例因子为 0.5 f 时，图像大小会减小如下：

$ ls -alh *.jpg
-rw-rw-r-- 1 ubuntu ubuntu 91K Nov 13 22:31 flower.jpg
-rw-rw-r-- 1 ubuntu ubuntu 23K Nov 17 02:46 output.jpg

使用 V100 测得的经过时间为0.04576 ms。它的时间可能会因 GPU 而异：

Rescale flower.jpg in 0.5 ratio.
CPU: 23.857 ms
GPU: 0.04576 ms
Done (generated output.jpg)

有关 NPP 用于图像处理的更多详细信息，请访问并查看链接的文档：on-demand.gputechconf.com/gtc/2014/presentations/HANDS-ON-LAB-S4793-image-processing-using-npp.pdf。

使用 NPP 进行信号处理

NPP 还提供了信号处理功能。与图像处理 API 的主要区别在于，它们不需要图像形状相关的信息。随着我们继续介绍 NPP 函数的基本用法，我们将了解如何从给定的数组中获取总和、最小/最大值、均值和 L2 归一化分布值。完整的代码是05_npp/statisticsNPP.cpp。

首先，让我们从所需的头文件开始：

#include <iostream>
#include <cuda_runtime.h>
#include <npp.h>

我们将使用随机生成的数字作为输入数据：

void GetData(float** buffer, size_t size)
{
    (*buffer) = (float*) new float[size];

    for (int i = 0; i < size; i++) {
        (*buffer)[i] = float(rand() % 0xFFFF) / RAND_MAX;
    }
}

在调用统计操作函数之前，我们需要一个临时内存空间进行操作。我们可以使用与操作相关的其他 NPP 函数来获取所需的大小，并创建一个公共工作空间内存空间：

int GetWorkspaceSize(int signalSize)
{
    int bufferSize, tempBufferSize;

    nppsSumGetBufferSize_32f(signalSize, &tempBufferSize);
    bufferSize = std::max(bufferSize, tempBufferSize);
    nppsMinGetBufferSize_32f(signalSize, &tempBufferSize);
    bufferSize = std::max(bufferSize, tempBufferSize);
    nppsMaxGetBufferSize_32f(signalSize, &tempBufferSize);
    bufferSize = std::max(bufferSize, tempBufferSize);
    nppsMeanGetBufferSize_32f(signalSize, &tempBufferSize);
    bufferSize = std::max(bufferSize, tempBufferSize);
    nppsNormDiffL2GetBufferSize_32f(signalSize, &tempBufferSize);
    bufferSize = std::max(bufferSize, tempBufferSize);

    return bufferSize;
}

让我们从main()函数开始。首先，我们将开始准备输入数据，并了解所需的工作空间内存空间。我们将准备两种输入数据类型，并使用 NPP 比较它们的差异：

GetData(&h_input1, buf_size);
GetData(&h_input2, buf_size);
workspace_size = GetWorkspaceSize(buf_size);

之后，我们将为输入/输出和工作空间分配 GPU 内存空间。我们还将按以下方式传输输入数据：

cudaMalloc((void **)&d_input1, buf_size * sizeof(float));
cudaMalloc((void **)&d_input2, buf_size * sizeof(float));
cudaMalloc((void **)&d_output, sizeof(float));
cudaMalloc((void **)&d_workspace, workspace_size * sizeof(Npp8u));

现在，让我们使用 NPP 函数进行一些简单的统计操作：

nppsSum_32f(d_input1, buf_size, d_output, d_workspace);
cudaMemcpy(&h_output, d_output, sizeof(float), cudaMemcpyDeviceToHost);
std::cout << "Sum: " << h_output << std::endl;

nppsMin_32f(d_input1, buf_size, d_output, d_workspace);
cudaMemcpy(&h_output, d_output, sizeof(float), cudaMemcpyDeviceToHost);
std::cout << "Min: " << h_output << std::endl;

nppsMax_32f(d_input1, buf_size, d_output, d_workspace);
cudaMemcpy(&h_output, d_output, sizeof(float), cudaMemcpyDeviceToHost);
std::cout << "Max: " << h_output << std::endl;

nppsMean_32f(d_input1, buf_size, d_output, d_workspace);
cudaMemcpy(&h_output, d_output, sizeof(float), cudaMemcpyDeviceToHost);
std::cout << "Mean: " << h_output << std::endl;

NPP 还提供了报告两个输入之间差异的函数，如下所示：

nppsNormDiff_L2_32f(d_input1, d_input2, buf_size, d_output, d_workspace); 
cudaMemcpy(&h_output, d_output, sizeof(float), cudaMemcpyDeviceToHost);
std::cout << "NormDiffL2: " << h_output << std::endl;

然后，我们终止使用的内存。之后，让我们用以下命令编译代码：

$ nvcc -run -m64 -std=c++11 -I/usr/local/cuda/samples/common/inc -gencode arch=compute_70,code=sm_70 -lnppc -lnppif -lnppisu -lnppig -lnppicom -lnpps -o statisticsNPP ./statisticsNPP.cpp

因此，我们得到的结果如下：

Sum: 0.00100016
Min: 1.30432e-06
Max: 3.04836e-05
Mean: 1.56275e-05
NormDiffL2: 9.46941e-05

NPP 的应用

在本节中，我们已经涵盖了图像处理中的滤波和信号处理中的统计操作。尽管我们尝试了简单的应用程序，但我们可能会发现 NPP 编程比内核实现要容易得多。因此，NPP 被应用于许多媒体转码滤镜、浴图像处理应用程序、计算机视觉或深度学习中的图像预处理等领域。

在 OpenCV 中编写 GPU 加速的代码

OpenCV 库在计算机视觉中非常受欢迎。它支持 GPU 编程，以便在计算机视觉领域的更高分辨率下获得更好的性能。在本节中，我们将介绍如何在 OpenGL 中使用 GPU。

CUDA-enabled OpenCV installation

要开始使用 CUDA 进行 OpenCV 编程，您需要使用启用 CUDA 功能的 OpenCV 库进行编译。按照以下步骤在 Ubuntu 中启用 OpenCV：

$ sudo apt-get install -y --no-install-recommends \
 cmake git libgtk2.0-dev pkg-config libavcodec-dev \
    libavformat-dev libswscale-dev \
 libatlas-base-dev gfortran libeigen3-dev \
 libgtkglext1 libgtkglext1-dev

如果您的系统可以使用 X 窗口（而不是服务器），请安装其他软件包以启用 GTK 对话框：

$ sudo apt-get install -y —no-install-recommends \
 Libgtkglext1 libgtkglext1-dev

下载源代码并使用以下命令解压它们。这是在撰写时测试的 OpenCV，这是最新的 OpenCV 版本：

# We are install OpenCV 4.1.1
OPENCV_VERSION=4.1.1
OPENCV_DIR=opencv

# Download OpenCV and contrib source codes
mkdir -p ${OPENCV_DIR}
wget -O ${OPENCV_DIR}/opencv-${OPENCV_VERSION}.tar.gz https://github.com/opencv/opencv/archive/${OPENCV_VERSION}.tar.gz
wget -O ${OPENCV_DIR}/opencv_contrib-${OPENCV_VERSION}.tar.gz https://github.com/opencv/opencv_contrib/archive/${OPENCV_VERSION}.tar.gz

# Untar the files
tar -C ${OPENCV_DIR} -xzf ${OPENCV_DIR}/opencv-${OPENCV_VERSION}.tar.gz
tar -C ${OPENCV_DIR} -xzf ${OPENCV_DIR}/opencv_contrib-${OPENCV_VERSION}.tar.gz

现在，让我们使用以下命令编译下载的源代码。如果需要，您可以添加其他选项。它的编译需要一些时间：

# Build the codes and install
cd ${OPENCV_DIR}/opencv-${OPENCV_VERSION}
mkdir build
cd build
cmake -D CMAKE_BUILD_TYPE=RELEASE \
 -D CMAKE_INSTALL_PREFIX=/usr/local \
 -D ENABLE_PRECOMPILED_HEADERS=OFF \
 -D OPENCV_GENERATE_PKGCONFIG=ON \
 -D WITH_CUDA=ON -D WITH_CUVID=OFF -D BUILD_opencv_cudacodec=OFF \
 -D ENABLE_FAST_MATH=1 \
 -D CUDA_FAST_MATH=1 \
 -D OPENCV_EXTRA_MODULES_PATH=../../opencv_contrib-${OPENCV_VERSION}/modules \
 -D WITH_CUBLAS=1 \
 -D PYTHON_DEFAULT_EXECUTABLE=`which python3` \
 -D INSTALL_PYTHON_EXAMPLES=ON \
 -D BUILD_EXAMPLES=ON ..
make -j$(nproc)
sudo make install -j$(nproc)

要确认安装，请使用以下命令：

$ pkg-config —cflags opencv4 -I/usr/local/include/opencv4/opencv -I/usr/local/include/opencv4

在 OpenCV 4 中，CUDA 相关的函数和类在 CUDA 命名空间中定义。例如，您可以使用此命令创建 CUDA 全局内存空间：

cv::cuda::GpuMat cuda_mem = cv::cuda::GpuMat(src.rows, src.cols, CV_8UC1);

然后，设备cuda_mem内存空间可以像正常的 CPU 内存类型（cv::Mat）一样处理。

实现 CUDA-enabled 模糊滤镜

现在，我们将实现一个小型的 GPU-enabled OpenCV 应用程序并比较其性能。让我们首先包括所需的头文件：

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

这是使用 OpenCV 的主机模糊滤镜实现：

void BlurHost(std::string filename)
{
    cv::Mat src = cv::imread(filename, 1);
    cv::Mat dst; 
    cv::TickMeter tm;

    tm.reset();
    tm.start();
    cv::bilateralFilter(src, dst, 10, 50, 50);
    tm.stop();

    std::cout << "CPU Time: " << tm.getTimeMilli() << " ms." << std::endl;
    cv::imwrite("result_host.jpg", dst);
}

这是 CUDA-enabled 模糊滤镜的实现：

void BlurCuda(std::string filename)
{
    cv::Mat src = cv::imread(filename, 1);
    cv::Mat dst;
    cv::cuda::GpuMat src_cuda = cv::cuda::GpuMat(src.rows, 
                                                 src.cols, CV_8UC1);
    cv::cuda::GpuMat dst_cuda = cv::cuda::GpuMat(src.rows, 
                                                 src.cols, CV_8UC1);
    cv::TickMeter tm;

    // warm-up
    cv::cuda::bilateralFilter(src_cuda, dst_cuda, 10, 50, 50);

    tm.reset();
    tm.start();
    src_cuda.upload(src);
    cv::cuda::bilateralFilter(src_cuda, dst_cuda, 10, 50, 50);
    dst_cuda.download(dst);
    tm.stop();

    std::cout << "GPU Time: " << tm.getTimeMilli() 
                  << " ms." << std::endl;
    cv::imwrite("result_cuda.jpg", dst);
}

此代码显示了bilateralFilter()操作如何与主机匹配，以及 CUDA 如何与 CUDA 命名空间匹配。对于 CUDA 内存操作，cv::cuda::GpuMat用于设备内存，并且设备内存提供upload()和download()成员函数，例如cudaMemcpy()。为了测量经过的时间，使用了cv::TickMeter。然后，main()调用了两种实现，如下所示：

  int main(int argc, char *argv[])
  {
      std::string filename("flower.JPG");

      BlurHost(filename);
      BlurCuda(filename);

      return 0;
  }

现在，让我们编译代码。我们应该在编译选项中使用pkg-config --cflag opencv包含 OpenCV 头文件和库。例如，编译选项可以写成这样：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc `pkg-config opencv4 --cflags --libs` -o blur ./blur.cpp

然后，输出结果如下：

CPU Time: 57.6544 ms.
GPU Time: 2.97784 ms.

执行时间可能会因系统和 GPU 而异。

启用多流处理

在 OpenCV 中，CUDA 流使用cv::cuda::Stream进行管理。使用这个，我们可以进行基于多流的管道化 GPU 操作：

我们知道，主机内存应该是固定的内存，以便进行异步数据传输：

Mat::setDefaultAllocator(cuda::HostMem::getAllocator(cuda::HostMem::PAGE_LOCKED));

然后，我们将创建多个流，如下所示：

const int num_stream = 4;
cuda::Stream stream[num_stream];

然后，我们加载源图像并根据加载的图像信息初始化 GPU 内存，如下所示：

Mat src = imread(filename, 1);
Mat dst;
cuda::GpuMat src_cuda[num_stream], dst_cuda[num_stream];
for (int i = 0; i < num_stream; i++)
    src_cuda[i] = cuda::GpuMat(src);

现在，我们将使用每个流将图像传输到 GPU，对图像进行模糊处理，然后将其传输回主机：

for (int i = 0; i < num_stream; i++) {
    src_cuda[i].upload(src, stream[i]);
    cuda::bilateralFilter(src_cuda[i], dst_cuda[i], 21, 150.f, 
                          150.f, BORDER_DEFAULT, stream[i]);
    dst_cuda[i].download(dst, stream[i]);
}

然后，我们必须同步主机和 GPU。为此，我们将使用cv::Stream.waitForCompletion()函数，在每个流完成数据传输到主机后进行同步：

for (int i = 0; i < num_stream; i++)
    stream[i].waitForCompletion();

为了与 CPU 性能进行比较，我们也调用cv::bilateralFilter()如下：

bilateralFilter(src, dst, 21, 150, 150);

其执行时间如下。GPU 执行时间是从多流执行循环到同步的测量时间的平均值：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc `pkg-config opencv4 --cflags --libs` -o blur ./blur.cpp
CPU Time: 84.8649 ms.
GPU Time: 1.60979 ms.

为了确认多流操作，我们可以对操作进行分析。以下截图显示了这一点：

对操作进行分析

默认流上的第一个操作是预热执行，接下来是四个多流操作。在这里，我们可以看到 GPU 操作是重叠的。因此，平均执行时间比单流执行时间短。

我们只介绍了 OpenCV 中的双边滤波。然而，许多 OpenCV 功能支持 CUDA 加速，因此您可以获得 GPU 计算的好处。其接口与 CPU 版本一致，因此您可以轻松将 CPU 版本迁移到 GPU。

作为入门级别，有一些来自 GTC 的有用材料：

建议您从 OpenCV 的参考指南开始：docs.opencv.org/4.1.1/d2/dbc/cuda_intro.html。

编写与 CUDA 兼容的 Python 代码

如今，许多人使用 Python 与 CUDA。它不仅作为二进制文件的粘合剂，还使我们能够直接编写 GPU 加速的代码。作为粘合语言，Python 可以调用来自 CUDA C/C++库的 API，使用pybind11 (github.com/pybind/pybind11) 或 SWIG (swig.org/)。但是，我们必须编写 CUDA C/C++代码并将其集成到 Python 应用程序中。

但是，有一些 Python 软件包—Numba、CuPy 和 PyCUDA—可以使用 Python 进行 GPU 编程。它们为 CUDA 内核提供了本机加速的 API 和包装器。换句话说，我们不必编写 C/C++代码并花费时间进行集成。Numba 提供了矢量化和 CUDA 即时（jit）编译器来加速其操作。它与 NumPy 兼容，因此您可以加速基于 NumPy 的数值计算代码。您还可以通过 jit 编译器在 Python 中编写灵活的 CUDA 代码。CuPy 也与 NumPy 兼容，并加速线性代数算法。它提供了 Pythonic 的可编程性和透明的自定义内核编程，如 Numba。PyCUDA 提供了 CUDA C/C++接口，因此您可以在 Python 代码中编写和使用 CUDA 内核函数。

Numba – 高性能 Python 编译器

Numba (numba.pydata.org/) 可以将 Python 函数翻译成在 GPU 上执行，而无需进行任何 C/C++编程。

在 Numba 中，您可以通过将 Numba 装饰器应用于目标函数轻松编写矢量化函数：

from numba import vectorize
@vectorize(["float32(float32, float32, float32)"], target='cuda')
def saxpy(scala, a, b):
return scala * a + b

如您所见，装饰器指定了参数和返回数据类型，目标指定了代码将在哪种架构上运行。有三种目标：

目标	描述	推荐的数据大小和操作
`cuda`	针对 NVIDIA GPU	大于 1MB，计算密集型操作
`parallel`	优化多核 CPU	小于 1MB，正常操作
`cpu`	优化单线程操作	小于 1KB，低计算密集型操作

如果您的函数不返回值，请使用@guvectorize，并将参数指定为向量。

Numba 的另一个用途是使用@cuda.jit装饰器。这使您能够编写类似以下的 CUDA 特定操作：

from numba import cuda

@cuda.jit
def matmul(d_c, d_a, d_b):
    x, y = cuda.grid(2)
    if (x < d_c.shape[0] and y < d_c.shape[1]):
        sum = 0
        for k in range(d_a.shape[1]):
            sum += d_a[x, k] * d_b[k, y]
        d_c[x, y] = sum

cuda.grid() 关键字提供了 CUDA 线程在网格级别的索引，因此您可以以 Python 的方式编写内核代码，例如 CUDA C/C++代码。调用 CUDA 内核函数可以按以下方式完成：

matmuldimGrid, dimBlock

现在，让我们安装这个软件包并尝试一些示例。

安装 Numba

要在 Python 代码中使用 Numba，您需要安装该软件包，并配置环境变量：

$ pip3 install numba
$ export NUMBAPRO_NVVM=/usr/local/cuda/nvvm/lib64/libnvvm.so
$ export NUMBAPRO_LIBDEVICE=/usr/local/cuda/nvvm/libdevice/

您需要将环境变量设置放在.bashrc或.zshrc的末尾，以便将来使用时更加方便。如果它们没有设置，Python 将返回此消息：

numba.cuda.cudadrv.error.NvvmSupportError: libNVVM cannot be found. Do `conda install cudatoolkit`:
library nvvm not found

使用带有@vectorize 装饰器的 Numba

我们将使用@vectorize装饰器测试一个简单的saxpy操作。这将把一个特定的函数转换为并行工作：

创建numba_saxpy.py。
导入numba、numpy和其他所需的软件包：

import numpy as np
from numba import vectorize
from timeit import default_timer as timer

使用@vectorize装饰器编写一个带有目标'cuda'的saxpy代码，以便在 CUDA 设备上工作：

@vectorize(["float32(float32, float32, float32)"], target='cuda')
def saxpy_cuda(scala, a, b):
    return scala * a + b

使用带有@vecotrize装饰器和目标为'parallel'的 Numba 编写 saxpy 代码，以在多核处理器（主机）上工作：

@vectorize(["float32(float32, float32, float32)"], target='parallel')
def saxpy_host(scala, a, b):
    return scala * a + b

编写一个操作代码，调用一些 NumPy 生成的输入数据的函数：

scala = 2.0
np.random.seed(2019)
print("size \t\t CUDA \t\t CPU")
for i in range(16,20):
    N = 1 << i
    a = np.random.rand(N).astype(np.float32)
    b = np.random.rand(N).astype(np.float32)
    c = np.zeros(N, dtype=np.float32)

    # warm-up
    c = saxpy_cuda(scala, a, b)

    # measuring execution time
    start = timer()
    c = saxpy_host(scala, a, b)
    elapsed_time_host= (timer() - start) * 1e3
    start = timer()
    c = saxpy_cuda(scala, a, b)
    elapsed_time_cuda = (timer() - start) * 1e3
    print("[%d]: \t%.3f ms\t %.3f ms" % (N, elapsed_time_cuda, elapsed_time_host))

这段代码报告了各种操作数大小的耗时：

size         CUDA        CPU
[65536]:   1.174 ms    0.199 ms
[131072]:  1.362 ms    0.201 ms
[262144]:  2.240 ms    0.284 ms
[524288]:  2.384 ms    0.337 ms

在这种情况下，CUDA 的性能比 CPU 慢，因为操作很简单，但数据传输开销很大。

使用带有@cuda.jit 装饰器的 Numba

我们还可以使用@cuda.jit装饰器编写复杂的操作，以在 GPU 上使用 Numba：

创建numba_matmul.py。
导入numpy、numba和其他所需的软件包：

import numpy as np
from numba import cuda
from timeit import default_timer as timer

使用@cuda.jit装饰器编写矩阵乘法代码：

@cuda.jit
def matmul(d_c, d_a, d_b):
    x, y = cuda.grid(2)
    if (x < d_c.shape[0] and y < d_c.shape[1]):
        sum = 0
        for k in range(d_a.shape[1]):
            sum += d_a[x, k] * d_b[k, y]
        d_c[x, y] = sum

在这段代码中，我们使用cuda.grid(dimension_size)来指定网格中的 CUDA 线程索引，因此，我们可以在 Python 中指定 CUDA 线程的索引。

创建a和b矩阵作为 NumPy 矩阵：

N = 8192
a = np.random.rand(N, N).astype(np.float32)
b = np.random.rand(N, N).astype(np.float32)

将 NumPy 生成的数据复制到设备：

d_a = cuda.to_device(a)
d_b = cuda.to_device(b)

创建将放置在 CUDA 设备内存中的c矩阵：

d_c = cuda.device_array((N, N))

调用矩阵乘法内核函数：

start = timer()
matmuldimGrid, dimBlock
elapsed_time_gpu = (timer() - start) * 1e3

将输出复制到主机：

c = d_c.copy_to_host()

将 CUDA 操作与主机进行比较：

# matrix multiplication (cpu)
start = timer()
c_host = np.matmul(a, b)
elapsed_time_cpu = (timer() - start) * 1e3

# print elapse times
print("Elapsed Time")
print("GPU: %.3f ms" % elapsed_time_gpu)
print("CPU: %.3f ms" % elapsed_time_cpu)

if (np.allclose(c_host, c)):
print("Done.")
else:
print("GPU and host results are mismatching.")

使用@cuda.jit装饰器和内置的cuda.grid()关键字，这个示例代码展示了在 Python 中实现 Numba 到矩阵乘法是多么简单。这段代码报告了设备和主机上的操作耗时：

Elapsed Time
GPU: 104.694 ms
CPU: 1539.005 ms
Done.

现在，让我们来介绍 CuPy，它可以在 CUDA 编程中实现更多的 Python 编程。

CuPy- GPU 加速 Python 矩阵库

CuPy (cupy.chainer.org) 使用 Python 实现线性代数加速，并通过使用 CUDA 库充分利用 GPU。它与 NumPy 兼容，并提供了愉快的 Python 编程体验。

让我们来介绍它的安装、基本用法和手动内核开发。

安装 CuPy

我们可以使用以下命令使用pip安装 CuPy。然后它还会安装cupy包和 CUDA 依赖项：

$ pip3 install cupy-cuda101    # CUDA 10.1
$ pip3 install cupy-cuda101    # CUDA 10.0
$ pip3 install cupy-cuda902    # CUDA 9.2

现在，让我们来介绍 CuPy 的基本用法。

CuPy 的基本用法

我们可以编写一个 saxpy 操作，如下所示：

>>> x = cp.arange(5).astype('f') 
>>> x 
array([0., 1., 2., 3., 4.], dtype=float32) 
>>> y = cp.arange(5).astype('f') 
>>> 0.5 * x + y 
array([0\. , 1.5, 3\. , 4.5, 6\. ], dtype=float32)

我们还可以使用matmul()函数进行矩阵乘法：

>>> x = cp.random.uniform(0, 1, (2, 4)).astype('float32') 
>>> y = cp.random.uniform(0, 1, (4, 2)).astype('float32') 
>>> cp.matmul(x, y)
array([[0.6514087, 0.826463 ], 
 [0.7826104, 0.2878886]], dtype=float32)

正如我们之前讨论的，CuPy 与 NumPy 兼容。基本上，之前的 CuPy 对象是 CuPy 的数组类型：

>>> x = cp.random.uniform(0, 1, (2, 4)).astype('float32') 
>>> type(x) 
<class 'cupy.core.core.ndarray'>

但是，我们可以使用cupy.asnumpy()函数将其转换为 NumPy 数组，如下所示：

type(cp.asnumpy(x))
<class 'numpy.ndarray'>

也可以使用cupy.asnumpy()函数进行反向操作。因此，我们可以基于这种兼容性进行以下操作：

>>> gpu = cp.random.uniform(0, 1, (2, 4)).astype('float32') 
>>> cpu = np.random.uniform(0, 1, (2, 4)).astype('float32')
>>> gpu + cp.asarray(cpu) 
array([[0.8649391 , 1.1412742 , 1.1280626 , 0.38262686],
 [0.44767308, 0.738155 , 0.8397665 , 1.5165564 ]], dtype=float32)
>>> cpu + cp.asnumpy(gpu) 
array([[0.8649391 , 1.1412742 , 1.1280626 , 0.38262686], 
 [0.44767308, 0.738155 , 0.8397665 , 1.5165564 ]], dtype=float32)

正如您所看到的，我们可以轻松切换目标计算过程，并且可以从每个平台的优势中受益。现在，让我们来介绍使用 CuPy 进行自定义内核实现。

实现自定义内核函数

CuPy 提供了三种自定义内核函数：elementwise、reduction 和 raw kernels。elementwise 内核有助于为每个元素进行自动索引。因此，我们只需编写一个元素的操作。reduction 内核执行减少操作，同时执行用户定义的操作。raw kernel 可以在 Python 代码上直接进行 CUDA C/C++内核编程，因此我们可以对其进行任何操作。在本节中，我们不会涵盖所有内容。但是，您可以从相关文档中了解更多信息-docs-cupy.chainer.org/en/stable/tutorial/kernel.html。

让我们讨论用户定义的逐元素内核实现。这是一个逐元素操作的示例：

>>> squared_diff = cp.ElementwiseKernel( 
...     'float32 x, float32 y', 
...     'float32 z', 
...     'z = (x - y) * (x - y)', 
...     'squared_diff')

然后，我们可以执行逐元素操作，而无需显式的索引操作：

>>> x = cp.random.uniform(0, 1, (2, 4)).astype('float32') 
>>> y = cp.random.uniform(0, 1, (2, 4)).astype('float32') 
>>> squared_diff(x, y) 
array([[0.54103416, 0.01342529, 0.01425287, 0.67101586], 
 [0.04841561, 0.09939388, 0.46790633, 0.00203693]], dtype=float32)
>>> squared_diff(x, 0.5) 
array([[0.23652133, 0.22603741, 0.08065639, 0.00647551], 
 [0.00029328, 0.07454127, 0.00666 , 0.18399356]], dtype=float32)

正如您在此代码中所看到的，CuPy 提供了一个高度 Pythonic 的接口，并且易于学习。还有许多内部例程，这些例程也与 NumPy 兼容—docs-cupy.chainer.org/en/stable/reference/routines.html。换句话说，当我们需要在 NumPy 中进行加速计算时，可以考虑使用 CuPy。

现在，我们将介绍 PyCUDA，它提供了直接的内核编程和隐式内存管理包装器。

PyCUDA – Pythonic access to CUDA API

PyCUDA（documen.tician.de/pycuda/）使我们能够在 Python 代码中编写 CUDA C/C++代码，并在不编译的情况下执行它们。通过这种方式，您可以编写 CUDA C/C++代码，这些代码是特定于 CUDA 的操作。但是，由于 PyCUDA 不会优化您的内核函数，因此您必须自行优化此代码。

这是使用 PyCUDA 生成的代码片段：

import pycuda.autoinit     # initialize CUDA devices
from pycuda import driver, compiler, gpuarray
from string import Template

kernel_code_template = Template("""
__global__ void matmul_kernel(float *d_C, float *d_A, float *d_B)
{
    int idx_x = blockIdx.x * blockDim.x + threadIdx.x;
    ...
}
""")

mod = compiler.SourceModule(kernel_code_template.substitute(MATRIX_SIZE=N))
matmul_kernel = mod.get_function("matmul_kernel")
matmul_kernel(driver.Out(C), driver.In(A), driver.In(B), block=(dimBlock, dimBlock, 1), grid=(dimGrid, dimGrid))

正如您在此代码中所看到的，我们可以使用相同的 Python 代码编写内核代码。我们还可以使用driver.In()和driver.Out()保留所需数据传输的标志。这表明 PyCUDA 在调用内核之前应传输数据。数据传输是自动的，我们也可以按以下方式传输数据：

d_A = driver.to_device(A) # cudaMemcpyHostToDevice
A = driver.from_device_like(d_A) # cudaMemcpyDeviceToHost

现在，让我们安装 PyCUDA 并尝试一些简单的示例。

安装 PyCUDA

要使用 PyCUDA，还需要安装该软件包。从网站（pypi.org/project/pycuda/）下载 PyCUDA 源文件。目前正在使用 2019.1.1 版本。

然后按照以下方式安装依赖项：

$ sudo apt-get install build-essential python-dev python-setuptools libboost-python-dev libboost-thread-dev

$ tar -xzf pycuda-2019.1.2.tar.gz
$ cd pycuda-2019.1.2
$ python3 ./configure.py --cuda-root=/usr/local/cuda --cudadrv-lib-dir=/usr/lib \
 --boost-inc-dir=/usr/include --boost-lib-dir=/usr/lib \
 --boost-python-libname=boost_python-py36 --boost-thread-libname=boost_thread
$ python3 setup.py build
$ sudo python3 setup.py install

如果要使用 Python 2，请跳过使用 Python 3 进行configure.py命令。配置命令可能会因您的 Python 版本而异。

使用 PyCUDA 进行矩阵乘法

我们可以使用 PyCUDA 以以下方式执行矩阵乘法：

创建一个pycuda_matmul.py文件。
导入所需的软件包如下：

import pycuda.autoinit
from pycuda import driver, compiler, gpuarray
import numpy as np
from string import Template
import timeit

编写 CUDA 内核函数代码：

kernel_code_template = Template("""
__global__ void matmul_kernel(float *d_C, float *d_A, float *d_B)
{
    int idx_x = blockIdx.x * blockDim.x + threadIdx.x;
    int idx_y = blockIdx.y * blockDim.y + threadIdx.y;

    float sum = 0.f;
    for (int e = 0; e < ${MATRIX_SIZE}; e++)
        sum += d_A[idx_y * ${MATRIX_SIZE} + e] * d_B[e * ${MATRIX_SIZE} + idx_x];
    d_C[idx_y * ${MATRIX_SIZE} + idx_x] = sum;
}
""")

使用 NumPy 生成输入/输出矩阵：

N = 8192
np.random.seed(2019)
A = np.random.rand(N, N).astype(np.float32)
B = np.random.rand(N, N).astype(np.float32)
C = np.zeros((N, N), dtype=np.float32)

编译内核代码：

mod = compiler.SourceModule( \
        kernel_code_template.substitute(MATRIX_SIZE=N))

从编译模块中获取内核函数：

matmul_kernel = mod.get_function("matmul_kernel")

使用从主机生成的输入数据创建设备内存：

d_A = gpuarray.to_gpu(A)
d_B = gpuarray.to_gpu(B)
d_C = gpuarray.zeros((N, N), dtype=np.float32)

配置网格和块维度：

dimBlock = 16
dimGrid = int((N + dimBlock - 1) / dimBlock)

准备获取 GPU 事件：

start = driver.Event()
stop = driver.Event()

调用内核函数：

print("Started GPU operation...")
start.record()

matmul_kernel(d_C, d_A, d_B, 
    block=(dimBlock, dimBlock, 1), 
    grid=(dimGrid, dimGrid))

stop.record()
stop.synchronize()
gpu_time = stop.time_since(start)
print("GPU Execution Time: %.3f ms" % (gpu_time))

从主机启动矩阵乘法，并将其与设备上的结果进行比较：

print("Started Host operation...")
start = timeit.default_timer()
c_host = np.matmul(A, B)
host_time = timeit.default_timer() - start

print("CPU Execution Time: %.3f ms" % (host_time * 1e3))

if (np.allclose(c_host, d_C.get())):
    print("Done.")
else:
    print("GPU and host results are mismatching.")

此代码还报告了设备和主机上的估计时间：

Started GPU operation...
GPU Execution Time: 657.547 ms
Started Host operation...
CPU Execution Time: 1531.133 ms
Done.

虽然 PyCUDA 公开了 CUDA C/C++内核代码，但由于性能不及 Numba 执行的操作，这表明需要手动优化内核。

Octave 和 R 中的零编码加速 NVBLAS

NVBLAS 是用于其他软件包（如 Octave 和 R）的 BLAS 操作的 CUDA 库。通过替换 OpenBLAS 执行的操作，Octave 或开发人员和数据科学家可以轻松享受 GPU 性能。在本章中，我们将介绍如何使用 NVBLAS 加速 Octave 和 R。

NVBLAS 是 cuBLAS 操作的动态库。cuBLAS 库是线性代数操作的 GPU 实现。它替换了 BLAS 库，因此我们可以轻松加速任何应用程序，而无需编码。让我们看看如何从 GEMM 示例代码中实现这一点。

配置

要在 Octave 和 R 中使用 NVBLAS，我们需要为 NVBLAS 提供一些工作环境变量。为此，让我们创建一个nvblas.conf文件，其中可以找到我们将使用 Octave 和 R 代码示例的目录。nvblas.conf文件可以编写如下：

NVBLAS_CPU_BLAS_LIB libopenblas.so
NVBLAS_LOGFILE nvblas.log
NVBLAS_GPU_LIST 0
NVBLAS_AUTOPIN_MEM_ENABLED

在这个文件中，我们可以看到 NVBLAS 需要了解 CPU 端的 BLAS 库。在本次会话中，我们将使用 OpenBLAS，因此需要使用以下命令在 Ubuntu 中安装它：

$ sudo apt-get install libopenblas-base libopenblas-dev

此外，我们可以通过为NVBLAS_GPU_LIST提供多个 GPU ID 来获得多 GPU 性能。本书提供了 GPU 执行结果，但如果您有多个 GPU，可以尝试提供多个 ID。

要在 Octave 和 R 中使用 NVBLAS，我们应该设置一个环境——LD_PRELOAD=libnvblas.so——与您的应用程序执行：

对于 Octave 代码，执行以下代码：

$ LD_PRELOAD=libnvblas.so octave sgemm.m

对于 R 脚本，执行以下命令：

$ LD_PRELOAD=libnvblas.so Rscript sgemm.R

当然，libnvblas.so文件应该可以从工作目录访问。它位于/usr/local/cuda/lib64/中。

NVBLAS 与存档软件包兼容。因此，使用以下命令安装的 Octave 和 R 软件包可以很好地与我们的测试配合使用：

$ sudo apt-get install octave # for octave installation
$ sudo apt-get install r-base # for R installation

现在，让我们尝试使用 Octave 和 R 语言来使用 NVBLAS。

加速 Octave 的计算

首先，我们将尝试使用 Octave 来使用 NVBLAS。完全实现的代码是08_nvblas/sgemm.m。实现如下：

for i = 1:5 
    N = 512*(2^i);
    A = single(rand(N,N));
    B = single(rand(N,N));

    start = clock();
    C = A * B;
    elapsedTime = etime(clock(), start);

    gFlops = 2*N*N*N/(elapsedTime * 1e+9);
    printf("Elapsed Time [%d]: %.3f ms, %.3f GFlops\n", N, elapsedTime, gFlops);
end

对于 GPU 操作，使用以下命令执行 Octave 脚本，并将性能与默认情况下的 GPU、NVBLAS 环境库和 CPU 进行比较：

$ LD_PRELOAD=libnvblas.so octave sgemm.m

然后，我们可以使用octave sgemm.m命令启动。输出结果如下：

CPU	GPU V100

Elapsed Time [1024]: 0.011 ms, 188.909 GFlops
Elapsed Time [2048]: 0.075 ms, 228.169 GFlops
Elapsed Time [4096]: 0.212 ms, 647.022 GFlops
Elapsed Time [8192]: 1.158 ms, 949.763 GFlops
Elapsed Time [16384]: 7.292 ms, 1206.241 GFlops

Elapsed Time [1024]: 0.010 ms, 208.346 GFlops
Elapsed Time [2048]: 0.024 ms, 721.731 GFlops
Elapsed Time [4096]: 0.094 ms, 1465.538 GFlops
Elapsed Time [8192]: 0.582 ms, 1889.193 GFlops
Elapsed Time [16384]: 4.472 ms, 1967.037 GFlops

可以看到，随着矩阵大小的增加，GPU 显示出更高的计算吞吐量。

加速 R 的计算

现在，我们将尝试使用 NVBLAS 来进行 R 语言的测试，以下是帮助步骤：

首先，让我们编写一个sgemm.R文件，执行点操作：

set.seed(2019)
for(i in seq(1:5)) {
    N = 512*(2^i)
    A = matrix(rnorm(N², mean=0, sd=1), nrow=N) 
    B = matrix(rnorm(N², mean=0, sd=1), nrow=N) 
    elapsedTime = system.time({C = A %*% B})[3]
    gFlops = 2*N*N*N/(elapsedTime * 1e+9);
    print(sprintf("Elapsed Time [%d]: %3.3f ms, %.3f GFlops", N, elapsedTime, gFlops))
}

使用以下命令执行 R 脚本，并比较性能：

$ LD_PRELOAD=libnvblas.so Rscript sgemm.R

样本代码多次操作，同时增加数据大小。以下表格显示了先前命令的输出：

CPU	GPU V100

Elapsed Time [1024]: 0.029 ms, 74.051 GFlops
Elapsed Time [2048]: 0.110 ms, 156.181 GFlops
Elapsed Time [4096]: 0.471 ms, 291.802 GFlops
Elapsed Time [8192]: 2.733 ms, 402.309 GFlops
Elapsed Time [16384]: 18.291 ms, 480.897 GFlops

Elapsed Time [1024]: 0.034 ms, 63.161 GFlops
Elapsed Time [2048]: 0.063 ms, 272.696 GFlops
Elapsed Time [4096]: 0.286 ms, 480.556 GFlops
Elapsed Time [8192]: 1.527 ms, 720.047 GFlops
Elapsed Time [16384]: 9.864 ms, 891.737 GFlops

从结果中，我们可以看到 CPU 和 GPU 之间的性能差距。此外，我们可以确定当样本大小增加时，GPU 的性能增益也会增加。

如果您对使用 GPU 加速的 R 感兴趣，请访问 NVIDIA 开发博客：devblogs.nvidia.com/accelerate-r-applications-cuda/

MATLAB 中的 CUDA 加速

MATLAB 是一个高效的、高级的数值分析工具，具有各种工具和函数。该工具从早期阶段就支持 CUDA，配备了并行计算工具箱。本节将向我们展示如何使用该工具生成 CUDA 代码。

要启用 GPU 加速，我们需要安装带有并行计算工具箱的 MATLAB。如果您已经有 MATLAB，请检查您的许可证是否包括并行计算工具箱。如果没有，您可以尝试 MATLAB 评估代码。从 MATLAB 的评估网站上，您可以下载任何类型的软件包，除了控制系统。大多数软件包都包含并行计算工具箱，所以您可以尝试这个。但是如果您不考虑使用 MATLAB，可以跳过这一部分。

当我们使用 MATLAB 代码在 GPU 上运行时，您需要使用 gpuArray 来创建设备内存。与Numba和PyCUDA将它们的主机数据发送到设备的方式相同，MATLAB 的gpuArray()创建设备内存并将给定的主机数据传输到设备上：

d_A = gpuArray(A);

本节假设您已经安装了 MATLAB 和并行计算工具箱。在本节中，我们将重点介绍实现示例代码，并比较主机和 GPU 的性能：

让我们编写一个在 CPU 上运行的host.m文件。代码如下：

N = 8192;
A = single(rand(N,N));
B = single(rand(N,N));

start = clock();
C = A * B; 
elapsedTime = etime(clock(), start);
gFlops = 2*N*N*N/(elapsedTime * 1e+9);
fprintf("Elapsed Time: %.3f ms, %.3f GFlops\n", elapsedTime, gFlops);

现在，让我们使用以下命令执行这两种实现。这是针对 MATLAB 的命令及其输出：

$ matlab -r "run('host.m'); exit;" -nodisplay
Elapsed Time: 6.421 ms, 171.243 Gflops

然后，让我们编写一个在 GPU 上运行的cuda.m文件。我们只需将gpuArray()应用于输入矩阵，如下所示：

N = 8192;
A = single(rand(N,N));
B = single(rand(N,N));

d_A = gpuArray(A);    'GPU memory allocation
d_B = gpuArray(B);    'GPU memory allocation

start = clock();
d_C = d_A * d_B;
elapsedTime = etime(clock(), start);
gFlops = 2*N*N*N/(elapsedTime * 1e+9);
fprintf("Elapsed Time: %.3f ms, %.3f GFlops\n", elapsedTime, gFlops);

这是 GPU 版本的执行代码和执行结果：

$ matlab -r "run('cuda.m'); exit;" -nodisplay
Elapsed Time: 0.179 ms, 6140.739 Gflops.

正如我们所看到的，GPU 相对于 CPU 表现出更高的性能。

MathWorks 提供了许多关于 MATLAB GPU 计算的示例。如果您想了解更多，请访问他们的网站：www.mathworks.com/examples/parallel-computing/category/gpu-computing。

总结

在这一章中，我们已经介绍了使用 CUDA 库和其他兼容语言的 CUDA 编程方法。我们还介绍了 cuBLAS 的基本用法及其混合精度操作特性。此外，我们还探讨了 cuRAND、cuFFT、NPP 和 OpenCV 库。由于这些库，我们可以轻松实现 GPU 应用程序，正如本章开头所讨论的那样。

我们已经使用其他与 CUDA 兼容的语言实现了一些 GPU 应用程序。首先，我们介绍了几个 Python 包，这些包使 Python 和 CUDA 可以互操作。它们提供了 Python 式的可编程性，并与其他 Python 特性兼容。然后，我们介绍了其他科学计算语言中的 CUDA 加速，例如 Octave、R 和 MATLAB。

现在，我们还有一种 GPU 编程方法要介绍——OpenACC。通过这种方法，我们可以使用诸如#pragma acc kernels之类的指令将原始的 C/C++和 Fortran 主机代码转换为在 GPU 上运行。我们将在下一章中介绍这一点。

第九章：使用 OpenACC 进行 GPU 编程

每个处理器架构都提供了不同的编写代码以在处理器上运行的方法。CUDA 也不例外；它也提供了不同的编码方法。近年来变得非常流行的一种方法是使用 OpenACC，它基本上是基于指令的编程。

OpenACC 基本上是一个将异构计算作为一等公民的标准。该标准基本上规定了有两种处理器，即主机和设备/加速器，这与 CUDA 编程模型所述的概念非常相似。

对于希望获得最佳性能的程序员来说，使用诸如 C、C++、Fortran 和 Python 等语言的 CUDA 编程是表达并行性的首选方式。编程语言要求程序员从头开始重新创建他们的顺序程序，同时保持他们的关键操作的串行和并行版本。并行编程语言中创建的并行程序往往只适用于非常少数的平台。

编译器指令将编程语言的灵活性与库的易用性相结合。程序员使用高级指令对代码进行注释，编译器可以使用这些指令来并行化代码，或者可以安全地忽略。这意味着带有编译器指令的代码可以编译为许多不同的并行平台，并且无需维护代码的串行和并行版本。此外，有时需要快速测试和原型化应用程序以在 GPU 上运行。一个这样的例子是将天气代码等代码库（拥有数百万行代码）转换为在 GPU 上运行；使用流行的语言将需要大量的工作。在这种情况下，OpenACC 成为一个合乎逻辑的选择。在 OpenACC 中，开发人员以指令的形式向编译器提供提示。编译器接受这些提示并生成特定于架构的加速器代码。

OpenACC 标准还为代码的开发人员提供了供应商中立性。带有 OpenACC 指令的单一源代码可以重新编译为不同的设备。例如，PGI 编译器目前支持 OpenACC 后端，如 Intel CPU 多核、NVIDIA GPU、Intel Xeon Phi 和 FPGA/ASIC 架构。这对于希望编写供应商中立代码的开发人员来说是一个非常有吸引力的提议。高性能计算中的关键应用程序，如 Vienna Ab-initio Simulation Package（VASP）（分子动力学/量子化学）、Weather Research and Forecasting（WRF）和 ANSYS Fluent（CFD）利用 OpenACC 编程模型来针对 NVIDIA GPU。

总结 OpenACC 的关键要点：

当异构计算被视为新的编程模型时，OpenACC 标准得以发展。
OpenACC 在各种加速器上提供性能可移植性。
OpenACC 并不是 CUDA 编程语言的替代品。当选择的处理器是 NVIDIA 时，OpenACC 编译器在后台生成 CUDA 代码。

近年来，OpenMP 标准也开始纳入异构计算 API。但迄今为止，还没有支持不同处理器架构的编译器，因此我们选择在本书中坚持使用 OpenACC。

本章将涵盖以下主题：

OpenACC 指令
OpenACC 中的异步编程
额外重要的指令和子句

技术要求

本章需要一台带有现代 NVIDIA GPU（Pascal 架构或更高）的 Linux/Windows PC。

如介绍中所述，OpenACC 是一个标准，这个标准由不同的编译器实现，如 GCC、PGI 和 CRAY 编译器。我们将在本章中使用的编译器是 PGI。PGI 编译器在 Fortran 社区中非常受欢迎，并且一直在实现 OpenACC 最新规范方面处于领先地位，并且提供了一个可以从 PGI 网站免费下载的社区版。好处是在社区版和付费版本的 PGI 编译器之间在功能上基本没有变化。在本章中，您需要下载 PGI 社区版。

本章的代码也可以在 GitHub 上找到：github.com/PacktPublishing/Learn-CUDA-Programming。

示例代码示例是使用 PGI 社区版的 19.4 版本开发和测试的。但建议您使用最新的 PGI 版本。

使用 OpenACC 在 GPU 上合并图像

为了理解 OpenACC 概念，我们选择了一个简单的计算机视觉算法来合并两个图像。在这段代码中，我们试图合并两个图像，如下所示：

前面的图像演示了一个计算机视觉算法，用于合并两个图像。

我们将在本章后面更多地讨论代码结构。首先，根据以下步骤配置环境：

准备您的 GPU 应用程序。例如，我们将使用一个用于合并两个图像的核算法。此代码可以在09_openacc/中找到。
使用pgc++编译器编译您的应用程序：

$ pgc++ -c -acc -ta=tesla:pinned scrImagePgmPpmPackage.cpp
$ pgc++ -c -acc -ta=tesla:pinned -Minfo=accel image_merging.cpp
$ pgc++ -o merging.out -acc -ta=tesla:pinned -Minfo=accel scrImagePgmPpmPackage.o image_merging.o
$ ./merging.out

前面的命令将创建一个名为blurring.out的二进制文件。正如您可能已经观察到的，我们正在使用pgc++编译器来编译我们的代码。此外，我们向我们的代码传递了一些参数。让我们更详细地了解它们：

-acc：此标志告诉编译器解析代码中提供的 OpenACC 指令。
-ta：代表应该为设备代码生成的目标架构。请注意，-ta=tesla表示我们的目标是 NVIDIA GPU。其他目标的一些示例包括-ta=multi-core，它将多核作为设备目标，-ta=radeaon，它将 AMD GPU 作为目标，还有一些其他目标。此外，我们可以添加特定于设备的标志；例如，我们为分配所有 CPU 内存作为固定（不可分页）的 GPU 添加了一个固定标志。
-Minfo：此选项告诉编译器为我们提供有关编译器采取的步骤的更多信息，使我们的代码并行化。通过说-Minfo-accel，我们要求编译器为我们提供与加速器区域相关的更多信息。我们可以将标志更改为-Minfo=all，以提供非加速器区域的详细信息。以下输出显示了向我们的代码添加Minfo标志的部分输出：

.... < More compiler output above>
merge_parallel_pragma(unsigned char *, unsigned char *, unsigned char *, long, long):
    30, Generating copyin(in1[:w*h])
    Generating copyout(out[:w*h])
    Generating copyin(in2[:w*h])
    Accelerator kernel generated
    Generating Tesla code
    30, #pragma acc loop gang /* blockIdx.x */
    32, #pragma acc loop vector(128) /* threadIdx.x */
    32, Loop is parallelizable
... < More compile output below >

要理解这个编译输出，我们需要了解 OpenACC pragma，我们将在下一节中进行。稍后我们将重新访问这个编译输出。可以使用pgc++ --help找到其他可用标志的更多详细信息。

运行二进制文件后的示例输出如下：

$ ./merging.out
Reading image width height and width [1536][2048]
Time taken for serial merge: 0.0028 seconds
Time taken for OpenACC merge(data+kernel): 0.0010 seconds
Time taken for OpenACC merge(kernel only) with Blocking: 0.0002 seconds
 Time taken for OpenACC merge(data _kernel) with blocking: 0.0014 seconds
Time taken for OpenACC merge (data+kernel)with Pipeline Async: 0.0008 seconds

前面的输出显示我们正在读取一个大小为 1536*2048 的图像。代码有一个串行实现和三个使用 OpenACC pragma 的并行实现。每个实现的时间在前面的输出中显示。最后一个使用 pipeline 方法的实现显示了最佳时间：0.0008 秒。我们将采取增量方法，并在接下来的部分详细介绍每个实现。

该算法的串行实现非常简单，如下面的代码片段所示：

void merge_serial(unsigned char *in1, unsigned char*in2, unsigned char *out, long w, long h)
{
    long x, y;
     for(y = 0; y < h; y++) {
         for(x = 0; x < w; x++) {
             out[y * w + x] = (in1[y * w + x]+in2[y * w + x])/2;
         }
     }
}

代码没有什么特别之处；基本上，它接受两个输入图像数据（in1和in2），执行平均操作以合并两个输入，最后存储输出。对于我们来说，关键的是循环是尴尬并行的，适合于 GPU 等架构。如上面的代码输出所示，串行实现花费了0.0028秒。请注意，计时可能会因运行代码的系统而略有不同。

在下一节中，我们将向您介绍 OpenACC 指令，以便将示例代码转换为在 GPU 上运行所需的指令。

OpenACC 指令

在本节中，我们将尝试理解 OpenACC pragma 的语法，并为合并操作实现基本的并行和数据指令。OpenACC pragma 的基本语法如下：

#pragma acc <directive> <clauses> 
!$acc parallel [clause [[,] clause]…]

上述命令解释如下：

在 C/C++中的#pragma被称为“编译器提示”。这些与程序员注释非常相似；但是，编译器实际上会读取我们的 pragma。如果编译器不理解 pragma，它可以忽略它，而不是抛出语法错误。
acc是我们 pragma 的一个补充。它指定这是一个 OpenACC pragma。任何非 OpenACC 编译器都会忽略此 pragma。
指令是 OpenACC 中的一个命令，它告诉编译器执行某些操作。目前，我们只会使用允许编译器并行化我们的代码的指令。
子句是对我们的指令的补充/修改。这些包括但不限于优化。

在本节中，我们将介绍三个指令：parallel，loop和data。我们将展示它们各自的用法，并最终将它们应用到我们的合并算法中。

并行和循环指令

并行指令是最直接的指令。它将标记代码的一个区域进行并行化（通常只涉及并行化一个for循环），如下面的代码所示：

#pragma acc parallel loop 
for (int i = 0; i < N; i++ ) {  
    //loop code 
}

我们还可以定义一个并行区域。并行区域可以有多个循环（尽管这通常不推荐！）。并行区域是指最外层花括号内的所有内容，如下面的代码片段所示：

#pragma acc parallel
{
    #pragma acc loop
    for (int i = 0; i < N; i++ )
    {
        < loop code >
    }
}

包含循环非常重要；否则，您将无法正确地并行化循环。并行指令告诉编译器冗余地并行化代码，如下所示：

循环指令明确告诉编译器我们希望并行化循环，如下面的屏幕截图所示：

循环指令有两个主要用途：

标记单个循环进行并行化
允许我们明确定义循环的优化/修改

我们将在本章后面讨论循环优化，以及 gang 和 vector；目前，我们将专注于并行化方面。循环指令要正常工作，必须包含在并行指令内：

#pragma acc parallel loop
for (int i = 0; i < N; i++ )
{
    //loop code 
}

使用并行指令时，必须包含循环指令才能使代码正常运行。我们还可以使用循环指令来并行化多维循环嵌套。在下面的代码片段中，我们看到了一个嵌套循环，并且我们明确为第二个循环提到了循环子句：

#pragma acc parallel loop
for (int i = 0; i < N; i++ )
{
    #pragma acc loop
    for( int j = 0; j < M; j++ )
    {
        //loop code
    }
}

请注意，在上面的代码片段中，我们没有在内部循环中再次放置并行子句，因为我们已经在从外部循环开始的范围中提到了它。

数据指令

OpenACC 并行模型规定我们有一个主机，运行我们的顺序代码（通常是 CPU）。然后我们有我们的设备，这是某种并行硬件。主机和设备通常（虽然并非总是）有单独的内存，程序员可以使用 OpenACC 在两个内存之间移动数据。

正如在第一章中讨论的，GPU 和 CPU 架构在根本上是不同的。GPU 作为吞吐量架构，具有大量计算单元和高速内存带宽。另一方面，CPU 是一种减少延迟的架构，具有大型缓存层次结构，并且提供大容量的主存储器。需要操作的任何数据都需要首先复制到 GPU 内存。（请注意，即使在统一内存的情况下，数据也会在后台以页面的形式由驱动程序复制。）

如下图所示，两种架构（CPU 和 GPU）之间的数据传输通过 I/O 总线进行：

在 OpenACC 中使用 GPU 作为目标架构的目标是仅将并行代码卸载到 GPU 上，而顺序代码将继续在 CPU 上运行。OpenACC 标准允许程序员通过使用 OpenACC 数据指令和数据子句 显式定义数据管理。数据子句允许程序员在主机和设备（或在我们的情况下，CPU 和 GPU）之间指定数据传输。

隐式数据管理：我们可以将数据传输留给编译器，如下例所示：

int *A = (int*) malloc(N * sizeof(int));

#pragma acc parallel loop
for( int i = 0; i < N; i++ )
{
    A[i] = 0;
}

在前面的代码中，编译器将理解需要从 GPU 复制A向量，并为开发人员生成隐式传输。

显式数据管理：最好使用显式数据传输来获得对传输更多控制，如下面的代码中使用复制数据子句所示：

int *a = (int*) malloc(N * sizeof(int));
#pragma acc parallel loop copy(a[0:N])
for( int i = 0; i < N; i++ )
{
     a[i] = 0;
}

在前面的代码片段中，我们使用了复制数据子句。下图解释了运行时到达复制数据指令时执行的步骤：

我们将通过合并代码的详细步骤来解释这些步骤，其中我们将应用数据子句。

其他可用的数据子句如下所列：

数据子句	描述	关键用法
`copy(list)`

在设备上分配内存
在进入区域时，从主机复制数据到设备
在退出区域时，将数据复制到主机

这是默认的输入数据结构，被修改后从函数返回
`copyin(list)`

在设备上分配内存
在进入区域时，从主机复制数据到设备

作为子例程的输入的向量
`copyout(list)`

在设备上分配内存
在退出区域时，将数据复制到主机

不覆盖输入数据结构的结果
`create(list)`

仅在设备上分配内存
不进行复制

临时数组

为了最大化性能，程序员应避免所有不必要的数据传输，因此显式内存管理优于隐式数据管理。

数组形状：数组形状是指定数组大小的方式。如果不指定形状，编译器将尝试假定大小。这在 Fortran 中效果很好，因为 Fortran 跟踪数组的大小；然而，在 C/C++中可能不起作用。数组形状也是从数组复制数据的唯一方式（例如，如果只需要复制数组的一半，这可能提高性能，减少不必要的复制），如下面的代码片段所示：

#pragma acc parallel loop copy(A[1:N-2])

这将复制A的所有元素，除了第一个和最后一个元素。

将并行、循环和数据指令应用于合并图像代码

现在让我们尝试将并行、循环和数据指令应用于合并顺序代码：

void merge_parallel_pragma(unsigned char *in1, unsigned char*in2,unsigned char *out, long w, long h)
{
    long x, y;
    #pragma acc parallel loop gang copyin(in1[:h*w],
                                          in2[:h*w]) 
                                          copyout(out[:h*w])
     for(y = 0; y < h; y++) {
        #pragma acc loop vector
        for(x = 0; x < w; x++) {
            out[y * w + x] = (in1[y * w + x]+in2[y * w + x])/2;
        }
    }
}

我们已经使用并行循环指令并行化了两个循环（高度：y和宽度：x）。此外，我们还明确地添加了数据子句来复制数据。请注意，由于in1和in2向量只是输入，它们是使用copyin()数据子句进行复制的。out向量是输出，使用copyout()数据子句进行复制。让我们试着理解这个函数的编译器输出：

merge_parallel_pragma(unsigned char *, unsigned char *, unsigned char *, long, long):
    30, Generating copyin(in1[:w*h])
        Generating copyout(out[:w*h])
        Generating copyin(in2[:w*h])
        Accelerator kernel generated
        Generating Tesla code
        30, #pragma acc loop gang /* blockIdx.x */
        32, #pragma acc loop vector(128) /* threadIdx.x */
32, Loop is parallelizable

前面的编译器输出显示，对于merge_parallel_pragma函数，编译器生成了以下操作：

在第 30 行，为in1和in2变量生成了copyin。在内核启动前将被复制到 GPU 的数组大小将是[0:w*h]。
在第 30 行，为out变量生成了copyout。在 GPU 内核启动后将被复制的数组大小将是[0:w*h]。
在第 30 和 32 行，生成了 Tesla 内核代码：
在第 30 行，外部循环使用了 gang 级并行化。
在第 32 行，内部循环使用了矢量级并行化

当代码在 V100 上运行时，整个内核所花费的时间为0.0010s。这基本上是串行代码的两倍快。这可能听起来并不令人印象深刻。原因是大部分时间花在了数据传输上，而不是内核计算。为了确认这一点，让我们使用nvprof：

$ nvprof ./merging.out
==26601== DoneProfiling application: ./merging.out
==26601== Profiling result:
Type Time(%) Time Calls Avg Min Max Name
GPU activities: 67.36% 609.41us 2 304.71us 286.34us 323.08us [CUDA memcpy HtoD]
27.63% 250.02us 1 250.02us 250.02us 250.02us [CUDA memcpy DtoH]
5.01% 45.344us 1 45.344us 45.344us 45.344us merge_parallel_pragma_30_gpu(unsigned char*, unsigned char*, unsigned char*, long, long)
...

正如您在前面的分析输出中所观察到的，94%的时间花在了数据传输上，而只有 5%的时间（45 微秒）花在了内核执行上。您可能会问：我怎么知道这是哪个内核？如果您仔细看 GPU 内核的名称，merge_parallel_pragma_30_gpu，PGI 编译器在merge_parallel_pragma函数的第 30 行生成了一个 CUDA 内核，这就是我们如何将其与在该行号放置的编译指示联系起来的方式。

所以我们知道问题在哪里，但解决方案呢？我们将使用的优化技术是 blocking 来隐藏这种延迟。我们将在接下来的章节中更多地介绍 blocking 技术，并使用异步子句来重叠这个传输。

OpenACC 中的异步编程

为了实现合并并行代码的更好性能，我们将利用一个叫做 blocking 的概念。Blocking 基本上意味着，我们可以创建数组的块，而不是一次性传输整个输入和输出数组，这些块可以并行传输和操作。以下图表演示了创建块并在内核执行时重叠数据传输：

前面的图表显示了不同的块被传输，这些块的内核执行可以独立于每个块。为了实现这一点，我们需要数据传输命令和内核调用被异步地触发和执行。为了实现 blocking，我们将在本节中引入更多的指令/子句：结构化/非结构化数据指令和async子句。我们将展示它们的每一个，并最终将它们应用到我们的基本 OpenACC 合并并行代码中。

结构化数据指令

OpenACC 数据指令允许程序员显式地管理设备上的数据（在我们的例子中是 GPU）。以下代码片段显示了标记结构化数据区域的示例：

< Initialize data on host (CPU) >
#pragma acc data < data clauses >
{
    //< Code >
}

设备内存分配发生在区域的开始，设备内存释放发生在区域的结束。此外，从主机到设备（CPU 到 GPU）的任何数据移动发生在区域的开始，从设备到主机（GPU 到 CPU）的任何数据移动发生在区域的结束。内存分配/释放和数据移动是由程序员包含的子句定义的。

包含多个计算区域：一个数据区域可以包含任意数量的并行/内核区域，如下例所示：

#pragma acc data copyin(A[0:N]) create(C[0:N])
{
    #pragma acc parallel loop
    for( int i = 0; i < N; i++ )
    {
        C[i] = A[i] + 10;
    }
    #pragma acc parallel loop
    for( int i = 0; i < N; i++ )
    {
        C[i] = C[i] / 10;
    }
}

非结构化数据指令

有两个非结构化数据指令：

进入数据：处理设备内存分配，并从主机复制到设备。您可以在进入数据中使用的两个子句是：
create：这将只执行设备内存分配。
copyin：这将执行分配以及从设备到设备的内存复制。
退出数据：处理设备内存释放，并从设备复制到主机。您可以在退出数据中使用的两个子句是：
delete：这将仅执行设备内存释放。
copyout：这将首先从设备复制内存到主机，然后执行设备内存释放。

非结构化数据指令不会将数据区域标记为您可以在代码中有多个进入数据和退出数据指令。最好将它们纯粹视为内存分配和释放。使用非结构化数据指令的最大优势是它们能够跨多个函数进行分支。您可以在一个函数中分配数据，并在另一个函数中释放它。我们可以看一个简单的例子：

#define N 1024
int* allocate(int size)
{
    int *ptr = (int*) malloc(size * sizeof(int));
    #pragma acc enter data create(ptr[0:size])
    return ptr;
} 
void deallocate(int *ptr)
{
    #pragma acc exit data delete(ptr)
    free(ptr);
}
int main()
{
    int *ptr = allocate(N);
    #pragma acc parallel loop
    for( int i = 0; i < N; i++ )
    {
        ptr[i] = 0;
    }
    deallocate(ptr);
}

上面的代码片段显示了分配发生在单独的allocate()函数中，删除发生在deallocate()中。您可以将相同的概念链接到 C++中构造函数的一部分enter data create和析构函数的一部分exit data delete。

OpenACC 中的异步编程

默认情况下，所有 OpenACC 调用都是同步的。这意味着，在每次数据传输或每次对 GPU 的内核调用之后，都会隐式添加同步。CPU 将等待直到 OpenACC 调用完成，然后开始执行下一条指令。为了使调用异步，我们可以在数据和并行指令中使用async子句，如下面的代码所示：

#pragma acc data copyin(a[:N]) async 
// performing copyin asynchronously 
#pragma acc parallel loop async 
//performing parallel loop asynchronously.

使用async的主要好处可以总结如下：

如果我们想要同时执行主机和设备代码，我们可以使用async启动我们的设备代码，而在执行时我们可以返回到主机继续不相关（非设备相关）的代码。
我们可以排队多个设备内核启动，以便它们连续执行，这在某些情况下可以减少启动设备内核的开销。
我们可以在主机和设备之间同时执行数据移动和设备计算。这是我们将应用于我们的代码的优化，并且是async的最常见用例。

在幕后，每当我们使用async子句时，我们都会向队列添加一些工作。提交给不同队列的工作可以异步执行，而在同一队列中的工作将顺序执行（一个接着一个）。当我们使用async时，我们可以指定队列号。如果未指定队列号，则将自动使用默认值。

将非结构化数据和异步指令应用于合并图像代码

现在让我们尝试将数据指令与async子句一起应用于合并并行代码：

void merge_async_pipelined(unsigned char *in1, unsigned char*in2,unsigned char *out, long w, long h)
{
    long x, y;
    #pragma acc enter data create(in1[:w*h], in2[:h*w], out[:w*h])
    const long numBlocks = 8;
    const long rowsPerBlock = (h+(numBlocks-1))/numBlocks;
    for(long block = 0; block < numBlocks; block++) {
        long lower = block*rowsPerBlock; // Compute Lower
        long upper = MIN(h, lower+rowsPerBlock); // Compute Upper
        #pragma acc update device(in1[lower*w:(upper-lower)*w],
                                  in2[lower*w:(upper-lower)*w]) 
                                  async(block%2)
        #pragma acc parallel loop present(in1,in2, out) async(block%2)
        for(y = lower; y < upper; y++) {
            #pragma acc loop
            for(x = 0; x < w; x++) {
                out[y * w + x] = (in1[y * w + x]+in2[y * w + x])/2;
            }
        }
        #pragma acc update self(out[lower*w:(upper-lower)*w]) 
                                async(block%2)
    }
#pragma acc wait
#pragma acc exit data delete(in1, in2, out)
}

我们已经使用了数据指令和async子句来实现阻塞概念。让我们分解整体实现，这将使其更容易理解：

进入数据区域：enter data create子句在 GPU 中为in1和in2变量以及out分配内存。
创建块：我们决定将图像分成八个块。这些块分布在行中。外部的for循环用于此目的添加了这个原因。
异步从主机传输数据到设备：acc update device基本上将数据从主机异步复制到设备，因为我们已经在其中添加了一个async子句。
异步启动并行循环：async子句被添加到并行子句中，以异步启动 GPU 内核。
异步从设备传输数据到主机：acc update self基本上是将数据从设备异步地复制到主机，因为我们已经在同一个地方添加了一个async子句。
等待：acc wait将确保 CPU 等待，直到所有 OpenACC 启动都完成，然后在所有队列中继续前进。
退出数据区域：acc exit data delete将删除在enter data子句中分配的in1和in2向量以及out。

让我们试着理解merge_async_pipelined函数的编译器输出：

merge_async_pipelined(unsigned char *, unsigned char *, 
                      unsigned char *, long, long):
     67, Generating enter data create(out[:h*w],in2[:h*w],in1[:h*w])
     74, Generating update device(in1[w*lower:w*(upper-lower)],
                                  in2[w*lower:w*(upper-lower)])
         Generating present(in1[:],out[:],in2[:])
         Accelerator kernel generated
         Generating Tesla code
         74, #pragma acc loop gang /* blockIdx.x */
         76, #pragma acc loop vector(128) /* threadIdx.x */
     76, Loop is parallelizable
     81, Generating update self(out[w*lower:w*(upper-lower)])
     84, Generating exit data delete(out[:1],in2[:1],in1[:1])

前面的编译器输出显示，对于merge_async_pipelined函数，编译器生成了以下操作：

在第 67 行，为in1、in2和out变量生成了data create区域。
在第 74 行，为in1和in2调用了update device，并且数据传输到设备被限制在上下界之间：in1[w*lower:w*(upper-lower)],in2[w*lower:w*(upper-lower)]。
在第 74 和 76 行，Tesla 内核代码已经生成。
在第 81 行，为out变量调用了update self，并且数据从设备传输被限制在上下界之间：out[w*lower:w*(upper-lower)]。
在第 84 行，数据区域结束，并调用delete来释放在 GPU 上分配的内存。

当代码在 V100 上运行时，整个内核所花费的时间为 0.0008 秒。为了更详细地了解这一点，让我们回到分析器。这次我们将利用 NVIDIA Visual Profiler 来可视化输出：

使用 NVIDIA Visual Profiler 输出

前面的屏幕截图显示了使用async和阻塞后的 Visual Profiler 输出。来自分析器窗口的关键消息如下：

我们看到有三个流被创建和使用。这是因为我们的代码使用了async(block%2)，这意味着我们请求了最大 2 个队列。第三个队列是默认队列，在管道执行期间不被使用。
我们看到主机到设备和设备到主机的传输也重叠了，因为 GPU 有两个直接内存访问（DMA）引擎，因此反向的数据传输也可以重叠。
我们还看到我们的内核执行与数据传输重叠。

到目前为止，我们已经看到了帮助我们将顺序代码转换为在 GPU 上运行的图像合并的关键指令。在下一节中，我们将向您介绍更多的子句，这些子句将帮助您进一步优化您的 OpenACC 代码。

其他重要的指令和子句

在本节中，我们将介绍其他重要的广泛使用的指令，可以应用到我们的合并算法中。

Gang/vector/worker

Gang/worker/vector 定义了我们可以在 OpenACC 中实现的各种并行级别。这种并行在并行化多维循环嵌套时非常有用。OpenACC 允许我们定义一个通用的 gang/worker/vector 模型，适用于各种硬件，但我们将更多地专注于 GPU 特定的实现。下图显示了 OpenACC 并行编程模型：

这个前面的图表代表了一个单一的 gang。当我们并行化我们的for循环时，循环迭代将会被均匀地分配给多个 gang。每个 gang 将包含一定数量的线程。这些线程被组织成块。一个 worker 是一行线程。

在前面的图中，有三个 worker，这意味着有三行线程。向量指的是每行有多长。所以在前面的图中，向量是八，因为每行有八个线程。在为 GPU 编程时，默认情况下会自动应用 gang 和 vector 并行。

由于 OpenACC 是一个开放标准并且面向多种硬件，它提供了通用构造。但是这个构造如何映射到特定的目标设备呢？答案很简单；这取决于架构和编译器，因此提供了性能可移植性。如果我们要映射当前 PGI 编译器如何将这个概念映射到 CUDA（NVIDIA GPU），那么它将如下所示：

OpenACC gang 映射到 CUDA 块。
worker 本质上映射到 CUDA 线程束。
OpenACC 向量映射到threadIdx.x和（X 维度）。
OpenACC worker 映射到threadIdx.y（Y 维度）。

再次强调，这是 PGI 编译器如何映射 OpenACC 构造的方式。其他编译器可能会以不同的方式进行映射。特别是对于 NVIDIA GPU，gang worker vector 将定义我们的 GPU 线程的组织。通过添加以下子句，开发人员可以告诉编译器在给定的循环上使用哪些并行级别：

gang: 标记用于 gang 并行的循环。
worker: 标记用于工作并行的循环。
vector: 标记用于向量并行的循环。

以下代码片段有三个循环，并且每个循环的并行性都已经明确定义：外循环为gang，中间循环为worker循环，最内层循环为vector循环：

#pragma acc parallel loop gang
for( i = 0; i < size; i++ )
    #pragma acc loop worker
    for( j = 0; j < size; j++ )
        #pragma acc loop vector
        for( k = 0; k < size; k++ )
          c[i][j] += a[i][k] * b[k][j];

调整 gangs、workers 和 vectors：编译器将为您选择一定数量的 gangs 和 workers 以及向量长度，但您可以使用以下子句进行更改：

num_gangs(N): 为并行区域生成N个 gangs
num_workers(M): 为并行区域生成M个 workers。
vector_length(Q): 为并行区域使用向量长度Q

在以下代码片段的示例中，我们将 gangs 的数量设置为2，workers 的数量设置为2，向量长度设置为32：

#pragma acc parallel num_gangs(2) \
  num_workers(2) vector_length(32)
{
  #pragma acc loop gang worker
  for(int x = 0; x < 4; x++){
    #pragma acc loop vector
    for(int y = 0; y < 32; y++){
      array[x][y]++;
    }
  }
}

在代码中设置 gangs 的数量很少是一个好主意——让编译器决定。大多数情况下，您可以通过调整向量长度有效地调整循环嵌套。此外，在 GPU 上很少使用 worker 循环。

托管内存

OpenACC 提供了一个选项，允许编译器处理内存管理。通过自己管理内存，我们将能够获得更好的性能；但是，允许编译器使用托管内存非常简单。我们不需要对我们的代码进行任何更改，就可以让托管内存正常工作。

为了使用托管内存，我们可以像这样将托管标志传递给pgc++编译器：

$ pgc++ -c -acc -ta=tesla:managed scrImagePgmPpmPackage.cpp
$ pgc++ -c -acc -ta=tesla:managed -Minfo=accel image_merging.cpp
$ pgc++ -o merging.out -acc -ta=tesla:managed -Minfo=accel scrImagePgmPpmPackage.o image_merging.o
$ ./blurring.out

添加了托管子句后，编译器基本上会忽略数据子句，并且托管内存用于在 CPU 和 GPU 之间传输数据。请注意，托管内存仅用于堆数据，而不是栈/静态数据。我们在上一章介绍的统一内存概念将保持不变。

内核指令

内核指令允许程序员退一步，完全依赖编译器。使用内核指令的一些示例代码如下：

#pragma acc kernels 
for (int i = 0; i < N; i++ ) 
{ 
    //< loop code > 
}

就像并行指令示例中一样，我们正在并行化一个循环。请记住，使用并行指令时，必须始终与循环指令配对；否则，代码将无法正确并行化。内核指令不遵循相同的规则；在一些编译器中，添加循环指令可能会限制编译器优化代码的能力。

内核指令是并行指令的完全相反。这意味着编译器做出了很多假设，甚至可能覆盖程序员并行化代码的决定。此外，默认情况下，编译器将尝试优化循环。编译器通常很擅长优化循环，并且有时甚至可以以程序员无法描述的方式优化循环。然而，通常程序员可以通过自己优化循环来获得更好的性能。

如果您遇到编译器拒绝并行化循环的情况，您可以覆盖编译器的决定。（但请记住，通过覆盖编译器的决定，您要对并行化代码造成的任何错误负责！）在这段代码中，我们使用独立子句来向编译器保证我们认为该循环是可以并行化的：

#pragma acc kernels loop independent
for (int i = 0; i < N; i++ )
{
    //< loop code >
}

Kernel 指令最明显的优势之一是它能够同时并行化许多循环。例如，在下面的代码段中，我们能够通过利用内核区域同时有效地并行化两个循环：

#pragma acc kernels
{
    for (int i = 0; i < N; i++ )
    {
        //< loop code >
    } 
... some other sequential code
    for (int j = 0; j < M; j++ )
    {
        //< loop code >
    }
}

Collapse 子句

collapse 子句允许我们将多维循环嵌套转换为单一维度循环。这个过程对于增加循环的整体长度（通常增加并行性）和通常有助于内存局部性。让我们看一下语法：

#pragma acc parallel loop collapse( 3 )
for(int i = 0; i < N; i++)
{
    for(int j = 0; j < M; j++)
    {
        for(int k = 0; k < Q; k++)
        {
            < loop code >
        }
    }
}

该代码将三维循环嵌套合并为单一维度循环。

Tile 子句

tile 子句允许我们将多维循环分解为瓦片或块。这通常对于增加某些代码的内存局部性很有用。让我们看一下语法：

#pragma acc parallel loop tile( 32, 32 )
for(int i = 0; i < N; i++)
{
    for(int j = 0; j < M; j++)
    {
        < loop code >
    }
}

前面的代码将我们的循环迭代分成 32 x 32 个瓦片（或块），然后并行执行这些块。

CUDA 互操作性

正如本章前面提到的，OpenACC 并不是 CUDA 语言的替代品；事实上，开发人员可以开始利用 OpenACC 将热点部分移植到 GPU 上。他们可以开始仅集成 CUDA 内核以用于最关键的功能。有几种方法可以将 OpenACC/CUDA 转换为可互操作的代码。我们将在本节中介绍其中一些。

DevicePtr 子句

这个子句可以用来映射使用cudaMalloc分配的 CUDA 设备指针，并将其传递给 OpenACC。以下代码片段展示了deviceptr子句的使用：

double *cuda_allocate(int size) {
    double *ptr;
    cudaMalloc((void**) &ptr, size * sizeof(double));
    return ptr;
}
int main() {
    double *cuda_ptr = cuda_allocate(100); 
    // Allocated on the device, but not the host!

    #pragma acc parallel loop deviceptr(cuda_ptr)
    for(int i = 0; i < 100; i++) {
        cuda_ptr[i] = 0.0;
    }
}

通常，OpenACC 运行时期望得到一个主机指针，然后将其转换为一些相关的设备指针。deviceptr子句是一种告诉 OpenACC 运行时一个给定指针不应该被转换，因为它已经是一个设备指针的方法。

Routine 指令

最后要讨论的话题是在 OpenACC 并行和内核区域内使用 CUDA 设备函数。这些函数是编译为由 GPU 内核或 OpenACC 区域调用的。为了在我们的 OpenACC 循环中使用 CUDA __device__函数，我们还可以使用 routine 指令：

//In CUDA code
extern "C" __device__
int cuda_func(int x) {
        return x*x;
}

//In OpenACC Code
#pragma acc routine seq
extern int cuda_func(int);

...

int main() {
    A = (int*) malloc(100 * sizeof(int));
    #pragma acc parallel loop copyout(A[:100])
    for(int i = 0; i < 100; i++) {
        A[i] = cuda_func(i);
    }
}

请注意，本章提供了一种实际利用 OpenACC 的方法，不涵盖整个标准 API。有关广泛的 API 信息，请参阅www.openacc.org/.

总结

在本章中，我们为您提供了一种利用 GPU 的替代方法。使用 OpenACC 的基于指令的编程方法对于传统应用程序非常受欢迎，对于新应用程序也提供了一种非常简单和可移植的方法。使用这种方法，您可以看到编译器变得更加先进。用户对指令的反馈已经被使用，通过利用指令可以为不同的架构生成最佳的并行代码。

我们介绍了提供指示/提示给编译器的并行指令。我们还利用数据指令来控制数据传输，而不是依赖于托管内存。通过使用异步子句，我们还尝试通过重叠内核和数据传输来优化我们的应用程序。我们探讨了将 OpenACC 构造映射到 CUDA 层次结构，以及 OpenACC 和 CUDA C/C++代码之间的互操作性。

在下一章中，我们将开始将我们对 CUDA 的知识应用于深度学习。

第十章：使用 CUDA 加速深度学习

深度学习是一种可以根据人工神经网络解释数据的机器学习方法。具体来说，我们提供机器可以理解的数据，并构建学习数据表示的神经网络模型。我们可以使用这种技术构建识别语音、从图像中分类对象、理解文本、翻译语言、转换数据域等模型。基本的神经网络包括全连接层（FCL）、卷积神经网络（CNN）和循环神经网络（RNN）。这些架构在数据分类、区域理解和顺序关系方面显示出强大的准确性。

深度学习需要大量计算，以便广泛应用。然而，通过使用 GPU 计算能力，我们可以显著减少训练时间，从而解决了这个问题。这是因为神经网络的基本架构是基于矩阵运算的，而 GPU 是一个针对此进行了优化的硬件平台。具体来说，深度学习的创新是通过 NVIDIA CUDA 加速来解决的，因为深度学习中的许多算法可以加速。

在本章中，我们将简要回顾神经网络操作，并讨论如何在 GPU 上加速这些操作。作为实践，我们将使用 cuDNN 和 cuBLAS CUDA 库实现一个卷积网络。cuDNN 库是 NVIDIA 的 CUDA 库，专门优化了深度学习操作。我们将在三个部分中介绍其实现。我们还将介绍 GPU 如何优化所需的操作。然后，我们将通过比较长短期记忆（LSTM）网络的性能来介绍使用 cuDNN 库的有效性。然后，我们将介绍使用NVIDIA 工具扩展（NVTX）进行深度学习的性能分析。这可以测量 GPU 上的网络操作，以便我们可以分析时间线上的操作并了解其性能。

在本章中，我们将涵盖以下主题：

使用 CUBLAS 加速全连接层
使用 cuDNN 的逐元素层
cuDNN/CUDA 中的 Softmax 和损失函数
使用 cuDNN 的卷积神经网络
使用 CUDA 的循环神经网络
深度学习框架的性能分析

技术要求

本章需要安装 cuDNN 库和 CUDA 工具包。我们还需要 CUDA 启用的 GPU。本章将介绍深度学习的基础知识和性能，因此不需要新的 GPU 功能。换句话说，如果您已经涵盖了前几章的大部分内容，您将拥有一个适当的 GPU 来使用。

要安装 cuDNN 库，您需要从developer.nvidia.com/cudnn下载软件包。您需要登录 NVIDIA 开发者网站才能访问下载页面。如果您还没有帐户，您需要注册一个 NVIDIA 开发者帐户。确保 cuDNN 与您安装的 CUDA 版本编译一致。

使用 cuBLAS 加速全连接层

全连接层是深度学习的基本架构。让我们回顾一下它的操作，并看看 CUDA 如何加速神经网络的前向和反向传播过程。然后，我们将把它们应用到 GPU 上。

神经网络操作

神经网络的基本操作是在输入数据和参数之间执行点操作。我们称之为感知。在深度学习中，神经网络以分层方式连接多个感知。我们称这些为前馈神经网络。以下图表显示了一个感知和基本神经网络：

感知器的基本操作是使用输入数据和适当的权重创建点积。然后，它使用激活函数进行非线性操作，例如 sigmoid 或整流线性单元（ReLU）。在前馈神经网络中，操作只是一个仿射变换，然后是激活函数的应用。一个向量将被馈送到神经网络作为输入，并与两层中每个节点之间的权重参数相乘。

为了训练神经网络，我们进行前向传播、损失计算和梯度反向传播，然后使用更新参数。让我们简要介绍一下它们。然后，我们将使用 cuBLAS 和其他 CUDA 操作来匹配每个步骤。

前向操作可以用以下方程表示：

这里，是给定输入向量的预测结果，是权重参数矩阵，是激活函数。正如我们所看到的，全连接层中的基本操作是矩阵运算。因此，我们需要对输入和激活函数实现矩阵乘法运算。因为我们进行分类任务，所以我们使用 softmax 函数来规范化输出，并在下一层获得概率分布结果。

为了获得真实值之间的损失，我们对标签应用 one-hot 编码，并通过从每个元素获得熵来获得交叉熵损失，如下所示：

我们可以通过每个交叉熵损失的总和来获得总损失值。然后，我们可以从前述方程中获得梯度。这看起来像一个复杂的操作，但可以简化如下：

现在，我们将梯度传播到前一层，这被称为反向传播。在这个任务中，我们使用链式法则来获得每个权重和偏差参数的梯度。然后，我们可以更新权重参数集和偏差。例如，我们可以通过以下方程获得权重和偏差的梯度：

我们可以通过以下方程获得梯度传播到前一层：

这里，是激活函数的梯度。因此，我们需要从第二层获得用于第一层。然后，可以通过以下方程获得第一层的权重和偏差的梯度：

现在，我们可以根据梯度下降规则更新权重和偏差，如下所示：

这里，是迭代步骤。

激活函数的梯度可能不同，其类型也可能不同。这个激活层的实现将在下一节中介绍。激活函数的导数可以用以下方程表示：

因此，神经网络操作是一组线性代数操作，并且可以使用 cuBLAS 库进行覆盖。实现的代码可以在01_ann中找到。我们将在实现全连接层、实现层操作和实现 softmax 层部分介绍这些实现细节。

神经网络层的设计

在编写代码之前，让我们来看看如何将操作打包成一个层配置：

首先，我们执行前向操作。
然后，我们执行反向操作。
然后我们从梯度中得到一个权重更新。
最后，输出层将获得损失。

这样，层可以配置如下：

它具有标准化的输入和输出，以及两种类型的输入，取决于工作流程。左侧数据路径将被命名为输入，而右侧将被命名为输出。数据分为两个阶段（前向和后向）。我们将使用 blob 来管理参数和输入/输出数据。blob 是跨层处理的数据的包装器，并帮助管理内存空间。我们将使用这种设计来简化网络的配置。每个层都将有每个 blob 的描述符和前向/后向处理操作。

现在，让我们创建一个层类，它将是所有层的基类。以下代码显示了class公共函数的堆叠。而且，你可以在01_ann/src/ directory的layer.h和layer.cu中找到它的实现。这不仅有前向和后向操作，还有权重更新控制和损失计算：

class Layer
{
public:
    Layer();
    ~Layer();

    std::string get_name() { return name_; }

    virtual Blob<float> *forward(Blob<float> *input) = 0;
    virtual Blob<float> *backward(Blob<float> *grad_input) = 0;

    virtual float get_loss(Blob<float> *target);
    virtual int   get_accuracy(Blob<float> *target);

    void set_cuda_context(CudaContext *context) { cuda_ = context; }

    /* weights update control */
    void freeze() { freeze_ = true; }
    void unfreeze() { freeze_ = false;}
    void set_load_pretrain() { load_pretrain_ = true; }
    void set_gradient_stop() { gradient_stop_ = true; }

为了支持这些操作，层类维护了几个 cuDNN 描述符、blob 指针和权重更新控制器。当我们涵盖网络实现时，详细的实现将会被涵盖：

protected:
    std::string name_;

    // Tensor descriptor for the input/output tensor
    cudnnTensorDescriptor_t input_desc_;
    cudnnTensorDescriptor_t output_desc_;
    // filter and bias descriptor for weights and biases
    cudnnFilterDescriptor_t filter_desc_;
    cudnnTensorDescriptor_t bias_desc_;

    // output memory
    Blob<float> *input_ = nullptr;       /* x */
    Blob<float> *output_ = nullptr;      /* y */
    Blob<float> *grad_input_ = nullptr;  /* dx */
    Blob<float> *grad_output_ = nullptr; /* dy */

    // master weights & bias
    bool freeze_ = false;               /* control parameter updates */
    Blob<float> *weights_ = nullptr;      /* w */
    Blob<float> *biases_  = nullptr;      /* b */
    Blob<float> *grad_weights_ = nullptr; /* dw */
    Blob<float> *grad_biases_  = nullptr; /* db */

    int batch_size_ = 0; // mini-batch size

    // cuda handle container
    CudaContext *cuda_ = nullptr;

    // initialize weights along with the input size
    void init_weight_bias(unsigned int seed = 0);
    void update_weights_biases(float learning_rate);

    // pretrain parameters
    bool load_pretrain_ = false;
    int load_parameter();
    int save_parameter();

    // gradient stop tagging
    bool gradient_stop_ = false;

    friend class Network;
}

这个层类将在其他部分的深度学习网络实现中使用。因此，它具有用于 cuDNN 操作的cudnnTensorDescriptor_t变量，以及get_loss()和get_accuracy()函数。

张量和参数容器

在我们的实现中，我们将使用一个名为Blob的数据容器。它的名称是从 Caffe 借来的。这使我们能够存储张量或网络参数以及其维度大小信息和内存点。我们将使用这个来连接每一层。这有助于每一层根据输入张量的大小信息初始化其权重。此外，每一层都可以根据Blob的信息验证其结果。

这个 blob 将需要神经网络中的维度大小信息，如下一行代码所示。然后，它的构造函数将根据大小信息创建一个主机端缓冲区：

Blob<T>(int n, int c, int h, int w)

Blob还可以处理主机和设备上的内存，并帮助我们访问这些内存。Blob具有以下内存访问辅助函数：

// get specified memory pointer
ftype *ptr() { return h_ptr_; }

// get cuda memory
ftype *cuda() 
{ 
    if (d_ptr_ == nullptr) 
        cudaMalloc((void**)&d_ptr_, sizeof(ftype) * len());
    return d_ptr_;
}

// transfer data between memory
ftype *to(DeviceType target) { 
    ftype *ptr = nullptr;
    if (target == host)
    {
        cudaMemcpy(h_ptr_, cuda(), sizeof(ftype) * len(), 
                   cudaMemcpyDeviceToHost);
        ptr = h_ptr_;
    }
    else // DeviceType::cuda
    {
        cudaMemcpy(cuda(), h_ptr_, sizeof(ftype) * len(), 
                   cudaMemcpyHostToDevice);
        ptr = d_ptr_;
    }
    return ptr;
}

正如我们之前讨论的，Blob可以存储张量，我们还需要提供张量形状信息，作为 cuDNN API 所需的描述符。因此，Blob可以使用以下代码创建和设置张量描述符：

/* Tensor Control */
bool is_tensor_ = false;
cudnnTensorDescriptor_t tensor_desc_;
cudnnTensorDescriptor_t tensor()
{
    if (is_tensor_)
        return tensor_desc_;

    cudnnCreateTensorDescriptor(&tensor_desc_);
    cudnnSetTensor4dDescriptor(tensor_desc_, 
                                CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT,
                                n_, c_, h_, w_);
    is_tensor_ = true;
    return tensor_desc_;
}

现在，让我们使用Blob来实现一个全连接层。

实现一个全连接层

在这一部分，我们将使用 cuBLAS 编写一个全连接网络。对于这个层，我们将创建一个从Layer类派生出来的Dense类。类构造函数将接收默认的层配置信息，如下所示：

Dense::Dense(std::string name, int output_size)
{
    name_ = name;
    output_size_ = output_size;
}

但这还不足以配置整个层。缺失的信息将从输入中提供，因为输入大小将由前一层确定。现在，让我们来看看前向传播。

实现前向传播

在前向传播中，我们可以将前向过程分为两个步骤，如下所示：

由于权重大小不必受批量大小的影响，我们只考虑输入权重和输出权重的数量。另一方面，数据馈送 blob，如输入和输出，受批量大小的影响。因此，我们的 GEMM 操作与过滤器和输入数据可以设计如下：

隐藏的输出将与偏置值相加。输入数据不仅限于数据加载器中的数据。当我们堆叠层时，上一层的输出将成为当前层的输入数据。前向操作可以实现如下：

Blob<float> *Dense::forward(Blob<float> *input) {
  .. { blob initialization } ..

  // output = weights^T * input (without biases)
  cublasSgemm(cuda_->cublas(),
        CUBLAS_OP_T, CUBLAS_OP_N, output_size_, 
        batch_size_, input_size_,
        &cuda_->one, weights_->cuda(), input_size_,
        input_->cuda(), input_size_,
        &cuda_->zero, output_->cuda(), output_size_);

  // output += biases * one_vec^T
  cublasSgemm(cuda_->cublas(), 
        CUBLAS_OP_N, CUBLAS_OP_N, output_size_, batch_size_, 1,
        &cuda_->one, biases_->cuda(), output_size_, one_vec, 1, 
        &cuda_->one, output_->cuda(), output_size_);
  return output_;
}

在第一次迭代中，每个层都需要初始化其权重和偏置。例如，这个Dense层可以初始化其权重、偏置和输出张量元素。我们可以将这个初始化任务分为两个阶段。第一个是权重和偏置，如下所示：

// initialize weights and biases
if (weights_ == nullptr)
{
    // setup parameter size information
    input_size_ = input->c() * input->h() * input->w();

    // initialize weight, bias, and output
    weights_ = new Blob<float>(1, 1, input_size_, output_size_);
    biases_ = new Blob<float>(1, 1, output_size_);
}

接下来的阶段是关于更新输入信息和初始化输出 blob。当它是新的或需要重新配置时，我们需要做以下工作。在这个任务中，我们还需要创建一个填满我们批量大小的向量。这将用于偏置的添加：

// initilaize input and output
if (input_ == nullptr || batch_size_ != input->n())
{
  input_ = input;
  batch_size_ = input->n();

  if (output_ == nullptr)
    output_ = new Blob<float>(batch_size_, output_size_);
  else
    output_->reset(batch_size_, output_size_);

  output_->tensor();

  if (d_one_vec != nullptr)
    cudaFree(d_one_vec);
  checkCudaErrors(cudaMalloc((void**)&d_one_vec, sizeof(float) * batch_size_));
  init_one_vec<<< (batch_size_+BLOCK_DIM_1D-1)/BLOCK_DIM_1D, BLOCK_DIM_1D >>>(d_one_vec, batch_size_);

  if (!freeze_)
    init_weight_bias();
}

这个初始化任务不仅触发了第一次迭代，还触发了批量大小的变化。在训练阶段不需要检查批量大小，但在测试阶段会很有用。这是因为训练和推断阶段的批量大小是不同的。在这种情况下，我们需要根据新的批量大小创建一个输出 blob。输出张量的大小是由通道大小确定的。以下代码创建了一个大小为（batch_size_，output_size_，1，1）的 blob：

output_ = new Blob<float>(batch_size_, output_size_);

这将创建扁平化张量。然后，我们将馈送这些张量，这要求它们在通道中对齐。这种对齐在 softmax 层中是特别需要的。我们将在 softmax 层的实现中进行讨论。

在这个阶段的另一个重要任务是初始化权重和偏置。在我们的实现中，我们将使用 ReLU 作为激活函数。我们将使用正常的初始化器（arxiv.org/abs/1502.01852）技术使网络可训练。根据前述论文的指导，所需的权重值可以用以下方程生成：

是来自上一层的输入数量。因此，我们可以在更新输入张量信息后初始化参数。此外，偏置值将被初始化为0。以下代码显示了这一实现：

void Layer::init_weight_bias(unsigned int seed)
{
    // Create random network
    std::random_device rd;
    std::mt19937 gen(seed == 0 ? rd() : static_cast<unsigned int>
                                        (seed));

    // He normal distribution
    float range = sqrt(6.f / input_->size());
    std::uniform_real_distribution<> dis(-range, range);

    for (int i = 0; i < weights_->len(); i++)
        weights_->ptr()[i] = static_cast<float>(dis(gen));
    for (int i = 0; i < biases_->len(); i++)
        biases_->ptr()[i] = 0.f;

    // copy initialized value to the device
    weights_->to(DeviceType::cuda);
    biases_->to(DeviceType::cuda);
}

现在，让我们来讨论反向传播。

实现反向传播

正如我们之前讨论的，来自下一层的梯度被传播到这一层。基于传播的梯度，我们需要获得权重、偏置和数据（输入梯度）的三个梯度。我们需要创建可以存储它们的 blob。它们的大小不取决于批量大小，所以我们只需要确保创建它们。以下代码显示了我们如何为此目的创建 blob：

if (grad_weights_ == nullptr) {
  grad_output_ = grad_output;
  grad_weights_ = new Blob<float>(weights_->shape());
  grad_biases_ = new Blob<float>(biases_->shape());
  grad_input_ = new Blob<float>(input_->shape());
}

在上述代码中，grad_output_表示从下一层传播的输出数据的梯度，grad_input_表示将传播到上一层的输入数据的梯度。因此，我们不需要创建grad_output_ blob。如果您觉得这些命名约定令人困惑，也许更容易理解grad_input_为，grad_input_为。

以下代码显示了我们如何实现这一点：

Blob<float> *Dense::backward(Blob<float> *grad_output) {
  .. { blob initialization } ..

  // db = (dy) * one_vec
  cublasSgemv(cuda_->cublas(),
    CUBLAS_OP_N,
    output_size_, batch_size_,
    &cuda_->one,
    grad_output_->cuda(), output_size_,
    one_vec, 1,
    &cuda_->zero,
    grad_biases_->cuda(), 1); 

  // dw = x * (dy)^T
  cublasSgemm(cuda_->cublas(),
    CUBLAS_OP_N, CUBLAS_OP_T,
    input_size_, output_size_, batch_size_,
    &cuda_->one,
    input_->cuda(), input_size_,
    grad_output_->cuda(), output_size_,
    &cuda_->zero,
    grad_weights_->cuda(), input_size_);

  // dx = W * dy
  if (!gradients_stop_)
    cublasSgemm(cuda_->cublas(),
      CUBLAS_OP_N, CUBLAS_OP_N,
      input_size_, batch_size_, output_size_,
      &cuda_->one,
      weights_->cuda(), input_size_,
      grad_output_->cuda(), output_size_,
      &cuda_->zero, 
      grad_input_->cuda(), input_size_);

  return grad_input_;
}

如果这一层是模型中的第一层，我们也可以跳过计算输入数据的梯度，因为我们不需要对其进行任何操作。

当我们想要更新权重时，将会更新权重和偏置值。在本节中，我们将使用随机梯度下降（SGD）来实现这一点。这个操作也可以在其他层中使用。在这里，我们将把这个函数放在Layer类中。权重更新也可以使用cublas函数来完成，如下所示：

void Layer::update_weights_biases(float learning_rate)
{
  float eps = -1.f * learning_rate;
  if (weights_ != nullptr && grad_weights_ != nullptr) {
    // w = w + eps * dw
    cublasSaxpy(cuda_->cublas(),
      weights_->len(),
      &eps,
      grad_weights_->cuda(), 1,
      weights_->cuda(), 1);
  }

  if (biases_ != nullptr && grad_biases_ != nullptr)
  {
    // b = b + eps * db
    cublasSaxpy(cuda_->cublas(),
      biases_->b(),
      &eps,
      grad_biases_->cuda(), 1,
      biases_->cuda(), 1);
  }
}

正如你所看到的，我们可以使用学习率更新权重和偏差。当然，你也可以改变eps操作以应用其他优化算法。

层终止

在 C/C++编程中，程序员应该覆盖如何在终止类实例时返回所使用的资源。根据我们的设计，如果层具有权重参数并且可以从梯度中更新它们，该层最多会创建六个 blob。以下代码显示了终止 blob 的层终止代码，这些 blob 是在内部创建的：

Layer::~Layer()
{
  if (output_ != nullptr) delete output_;
  if (grad_input_ != nullptr) delete grad_input_;

  if (weights_ != nullptr) delete weights_;
  if (biases_ != nullptr) delete biases_;
  if (grad_weights_ != nullptr) delete grad_weights_;
  if (grad_biases_ != nullptr) delete grad_biases_;
}

输入 blob 或张量描述符将由其他层或 blob 终止处理。层类是其他层的基类。因此，我们可以专注于终止自定义创建的资源，因为当我们终止任何派生层时，这个终止代码将一起被调用。

尽管我们已经设计了网络和层，但我们还应该开发一些额外的层来完成网络。例如，我们没有实现激活、softmax 和损失计算层。我们将在接下来的部分中介绍这些层。

使用 cuDNN 的激活层

神经网络层中有许多逐元素操作。激活函数是这些操作之一。cuDNN 库提供了六种激活函数：sigmoid、ReLU、tanh、clipped ReLU、ELU 和 identity。在 cuDNN 库中，cudnnActivationForward()执行前向操作，cudnnActivationBackward()执行后向操作。

让我们看一下cuddnnActivationForward()函数的接口，如下所示：

cudnnStatus_t cudnnActivationForward( cudnnHandle_t handle,
    cudnnActivationDescriptor_t activationDesc,
    const void *alpha, const cudnnTensorDescriptor_t xDesc, 
    const void *x, const void *beta,  
    const cudnnTensorDescriptor_t yDesc, void *y)

使用cudnnActivationDescriptor_t，我们可以确定激活函数的类型。Alpha 和 beta 是标量值，用于确定要添加的输入速率。xDesc和yDesc保存张量的形状信息。它们可以使用cudnnCreateTensorDescriptor()创建。

当你看cudnnActivationBackward()函数时，dy是来自下一层的梯度输入，dx是输出到上一层的梯度。在这种情况下，y变成了输入。这样，dyDesc提供了梯度输入形状信息，而dxDesc提供了梯度输出形状信息：

cudnnStatus_t cudnnActivationBackward( cudnnHandle_t handle,
    cudnnActivationDescriptor_t activationDesc,
    const void *alpha, const cudnnTensorDescriptor_t yDesc,  
    const void *y,
    const cudnnTensorDescriptor_t dyDesc, const void *dy,
    const cudnnTensorDescriptor_t xDesc,  const void *x,
    const void *beta,  const cudnnTensorDescriptor_t dxDesc, void *dx)

一般来说，我们可以期望层之间的张量形状不会改变。因此，我们可以对x和dx使用相同的张量描述符。这与使用y和dy是一样的。

现在，让我们使用 cuDNN API 实现启用 cuDNN 的激活函数。要使用 cuDNN API，我们需要提供一个张量描述符来指定输入和输出张量的维度给 cuDNN 函数。我们还需要指定激活操作。

层配置和初始化

虽然我们的示例实现没有使用层接口，但我们需要将我们的示例集成到层接口中。在我们的层设计中，激活层可以这样实现：

class Activation: public Layer
{
public:
  Activation(std::string name, cudnnActivationMode_t mode, 
             float coef = 0.f);
  ~Activation();

  Blob<float> *forward(Blob<float> *input);
  Blob<float> *backward(Blob<float> *grad_input);

private:
  cudnnActivationDescriptor_t act_desc_;
  cudnnActivationMode_t mode_;
  float coef_;
};

在初始化步骤中，我们需要创建几个张量描述符和一个激活描述符。cuDNN 库要求开发人员提供与 API 对应的张量大小或任何其他操作句柄：

Activation::Activation(std::string name, cudnnActivationMode_t mode, float coef)
{
  name_ = name;
  mode_ = mode;
  coef_ = coef;

  cudnnCreateActivationDescriptor(&act_desc_);
  cudnnSetActivationDescriptor(act_desc_, mode, CUDNN_PROPAGATE_NAN, coef);
}

在 cuDNN 中，我们使用激活描述符来指定激活函数操作。我们使用cudnnSetActivationDescriptor()函数来实现这一点。然后，它可以确定cudnnActivationForward/Backward()函数的操作。我们将在下一节中介绍这一点。然而，在这之前，我们需要实现类析构函数，以便它销毁激活描述符，如下所示：

cudnnDestroyActivationDescriptor(activation_desc);

现在，让我们介绍激活层的前向和后向操作。

实现层操作

这也被称为警告操作。这个层不需要我们处理权重和偏差，因此比密集层更容易实现。

实现前向传播

在第一次迭代中，我们需要初始化输入描述符、输出描述符和输出 blob。当批处理大小改变时，我们将更新输出 blob。然而，我们不需要初始化权重和偏差，因为它们没有。以下代码显示了它的实现：

if (input_ == nullptr || batch_size_ != input->n())
{
  input_ = input;
  input_desc_ = input->tensor();
  batch_size_ = input->n();

  if (output_ == nullptr)
    output_ = new Blob<float>(input->shape());
  else
    output_->reset(input->shape());

  output_desc_ = output_->tensor();
}

初始化后，我们使用 cuDNN 中的cudnnActivationForward()函数进行激活过程，如下所示：

cudnnActivationForward(cudnnHandle, act_desc_, 
    &one, input_desc_, d_input, &zero, output_desc_, d_output);

这个激活函数的操作是在我们初始化这个层时确定的，正如我们之前讨论的。

实现反向传播

下一步是实现反向传播。我们将重用我们已经拥有的输入/输出张量描述符。现在，我们必须初始化我们希望反向传播的梯度：

if (grad_input_ != grad_output_)
{
  grad_output_ = grad_output;
  grad_input_ = new Blob<float>(input_->shape());
  grad_input_->reset(input_->shape()); 
}

初始化后，我们可以调用cudnnActivationBackward()函数，如下所示：

cudnnActivationBackward(cudnnHandle, activation_desc, 
    &one, output_desc_, output_->cuda(), output_desc_, 
    d_grad_output, input_desc_, input_->cuda(),
    &zero, input_desc_, grad_input_->cuda());

请注意，我们重用了在前向传递中创建的输入张量描述符和输出张量描述符。我们之所以能够这样做，是因为激活操作不会改变张量的大小。我们可以通过在激活反向传播中使用 cuDNN API 来简化我们的实现。

cudnnActivationBackward()函数的输出是d_grad_input。正如我们在前一节中描述的，这个梯度将传递给下一层。

现在，我们将实现 softmax 层，并将我们的层实现集成为一个网络。然后，我们将讨论图像分类任务中全连接层的准确性。

cuDNN/CUDA 中的 softmax 和损失函数

对于 MNIST 数据集分类，我们将使用 softmax 分类器。softmax 函数对输入进行归一化，并生成概率的概率分布。softmax 操作可以表示如下：

cuDNN 的 softmax 前向函数支持此操作，以及通道和所有实例。之前，我们将密集层的输出与通道对齐。因此，我们将沿着通道应用 softmax 操作。

为了确认我们的训练有效完成，我们需要计算损失函数。由于 softmax 损失函数用于获取跨概率的损失，所以 softmax 损失函数被称为交叉熵损失。损失函数如下：

我们需要获得这个 softmax 损失的梯度以更新神经网络。幸运的是，softmax 损失的梯度在求导后很简单，如下所示：

对于前向操作，我们将使用 cuDNN 函数来获取 softmax 的输出。为了获得梯度，拥有自定义操作更直观和简单。

实现 softmax 层

现在，让我们看看如何使用 cuDNN 和 CUDA 代码来实现 softmax 层。

实现前向传播

我们可以使用 cuDNN 库中的cudnnSoftmaxForward()来获得 softmax 成本函数的输出：

cudnnSoftmaxForward(cudnnHandle, CUDNN_SOFTMAX_ACCURATE, 
      CUDNN_SOFTMAX_MODE_CHANNEL,
      &one,  input_desc,  d_input, &zero, output_desc, d_output);

在这种情况下使用的最重要的参数设置之一是CUDNN_SOFTMAX_MODE_CHANNEL。此选项使得在输入张量描述符信息后面进行通道级别的 softmax 操作。通过这样做，我们可以提供已经通过密集层的小批量输入按通道对齐的张量。

实现反向传播

softmax 层的反向传递与其他层的实现不同。这个操作将输入数据的标签作为输入，并获得适当的梯度。正如我们之前讨论的，softmax 损失的梯度可以使用以下方程获得：

我们可以使用cublasSaxpy()来实现这个操作，如下所示：

// set grad_input_ as predict
cudaMemcpyAsync(grad_input_->cuda(), output_->cuda(), 
                output_->buf_size(), cudaMemcpyDeviceToDevice));
// set grad_input_ = predict - target 
cublasSaxpy(cuda_->cublas(), target->len(), &cuda_->minus_one,
            target->cuda(), 1, grad_input_->cuda(), 1));

在前面的代码中，目标 blob 包含了 one-hot 编码的目标向量，因此将负目标向量添加到预测值中会产生适当的梯度。之后，我们需要在传播到前一层之前对批次梯度进行归一化，如下所示：

int grad_output_size = target->n() * target->c() * target->h() * target->w();
float scale = 1.0f / static_cast<float>(target->n());
cublasSscal(cuda_->cublas(), grad_output_size, &scale, grad_input_->cuda(), 1);

由于这引入了加权和的均值，我们可以期望每个批次的梯度被归一化。

实现损失函数

计算 softmax 的损失值是可选的。这意味着它的值在训练和推断中不被考虑。然而，我们可以将其用作训练的指标。

如我们之前讨论的，softmax 损失函数应该实现以下方程：

我们可以通过一个核函数从每个样本的输出中获得损失并累积它们，如下所示：

__global__ void
softmax_loss_kernel(float *reduced_loss, float *predict, 
                    float *target, int size)
{
  int batch_idx = blockDim.x * blockIdx.x + threadIdx.x;

  extern __shared__ float s_data[];
  float loss = 0.f;

  // each thread calculate entropy for each data 
  // and accumulate to shared memory
  if (batch_idx > 0)
    return;

  for (int c = 0; c < num_outputs; c++)
    loss += target[batch_idx * num_outputs + c] * \
                logf(predict[batch_idx * num_outputs + c]);
                workspace[batch_idx] = -loss;

  // Then, we do reduction the result to calculate loss 
  // Using 1 thread block
  if (blockIdx.x > 0) return;

  // Cumulate workspace data
  s_data[threadIdx.x] = 0.f;
  for (int i = 0; i < batch_size; i += blockDim.x)
    s_data[threadIdx.x] += workspace[threadIdx.x + i];

  __syncthreads();

  // Reduction
  for (unsigned int stride = blockDim.x / 2; stride > 0; stride >>= 1)
  {
    if (threadIdx.x + stride < batch_size)
      s_data[threadIdx.x] += s_data[threadIdx.x + stride];
    __syncthreads();
  }

  if (threadIdx.x == 0)
    reduced_loss[blockIdx.x] = s_data[0];
}

这个操作使用并行归约，在第三章 CUDA 线程编程中介绍过，用于获取一个批次中的累积损失值。由于我们只会使用这个减少的损失值来确认训练，所以我们只会监视它的输出而不是取平均值。

现在，让我们将我们实现的所有层与 MNIST 数据集加载器集成在一起。

MNIST 数据加载器

整个过程中一个重要的部分是为特定数据集创建一个数据加载器。在这个实验室中，我们将使用包含 60,000 个样本的 MNIST 数据集。在初始化时，我们告诉数据加载器它应该加载训练集还是测试集。之后，数据加载器将加载数据集中的一些魔术数字，以及所有样本和它们的标签。加载的数据将被存储在向量中，并使用相同的随机种子进行洗牌。由于数据加载器构建和洗牌样本向量，训练循环或测试循环可能会在每次迭代时获得随机化的输入数据。完整的实现代码可以在本书的 GitHub 存储库中的src/mnist.cpp文件中找到。

管理和创建模型

当我们有多个层时，我们需要一个可以管理这些层的对象，进行神经网络操作，即前向/后向传播和权重更新。在这个实验室中，我们将有一个层的数组，并迭代数组进行前向处理。例如，前向操作可以用以下代码执行：

Blob<float> *Network::forward(Blob<float> *input) {
  output_ = input;
  for (auto layer : layers_)
    output_ = layer->forward(output_);

  return output_;
}

反向传播也可以通过以相反顺序迭代数组来完成：

void Network::backward(Blob<float> *target) {
  Blob<float> *gradient = target;
  // back propagation.. update weights internally.....
  for (auto layer = layers_.rbegin(); layer != layers_.rend(); layer++) {
    // getting back propagation status with gradient size
    gradient = (*layer)->backward(gradient);
  }
}

如您所见，我们在向量中管理层，并具有每个层的操作。将新层添加到网络中甚至更简单，如下面的代码所示：

void Network::add_layer(Layer *layer) {
  layers_.push_back(layer);
}

通过使用Network类，我们可以使用各种模型管理函数，如参数更新，层注册，层初始化等。此外，我们可以构建一个像现代深度学习框架一样的神经网络。例如，我们可以创建一个模型如下：

// step 1\. loading dataset
MNIST data_loader = MNIST("./dataset");
// create training dataset loader and shuffling the data
data_loader.train(batch_size, true);  

// step 2\. model initialization
Network model;
model.add_layer(new Dense("dense1", 500));  // 1st layer
model.add_layer(new Dense("dense2", 10));   // 2nd layer
model.cuda();     // set cuda context for each layer

我们还可以有以下训练循环：

// get data sample's shared buffer
Blob<float> *train_data   = data_loader.get_data();   
// get target's shared buffer
Blob<float> *train_target = data_loader.get_target(); 
// load data and targets with the batch size
data_loader.get_batch();    
tp_count = 0;  step = 0;
while (step < num_steps)
{
  // transfer loaded data to the GPU
  train_data->to(cuda);
  train_target->to(cuda);

  model.forward(train_data);    // forward
  model.backward(train_target); // backward
  learning_rate *= 1.f / (1.f + lr_decay * step);
  model.update(learning_rate);  // update

  step = data_loader.next(true); // load next data

  ... monitoring logic ...
}

对于测试阶段，我们为测试数据集创建另一个数据集加载器，并只进行前向传播的迭代。以下代码显示了它的实现：

test_data_loader.test(batch_size_test);                   // create test dataset loader
Blob<float> *test_data = test_data_loader.get_data();     // get sample data shared buffer
Blob<float> *test_target = test_data_loader.get_target(); // get target shared buffer
test_data_loader.get_batch();    // load samples and targets with the batch size
tp_count = 0; step = 0;
while (step < num_steps_test) {
  // transfer loaded data to the GPU
  test_data->to(cuda);
  test_target->to(cuda);

  model.forward(test_data);  // forward
  tp_count += model.get_accuracy(test_target);

  step = test_data_loader.next(); // load next data
}
float accuracy = 100.f * tp_count / num_steps_test / batch_size_test;

在测试阶段，我们将在完成对测试数据集中所有样本的测试后获得准确率。现在，我们需要在测试循环之后获得准确率。

使用 MNIST 数据集进行网络训练

现在，让我们运行我们实现的代码并查看其结果。对于训练阶段，我们将迭代 2,400 步，批量大小为 256。MNIST 数据集在训练集中有 60,000 个样本。2,400 步意味着我们将进行大约 10 个 epochs 的迭代。样本代码可以用以下命令编译：

$ nvcc -run -m64 -std=c++11 -I/usr/local/cuda/samples/common/inc -gencode arch=compute_70,code=sm_70 -lcublas -lcudnn -lnvToolsExt -o train ./train.cpp ./src/layer.cu ./src/loss.cu ./src/mnist.cpp ./src/network.cpp

以下截图显示了我们实现的训练和测试输出：

在训练迭代中，网络从训练数据集中获得了 92%的准确率。然而，测试准确率只有 77%，这与训练结果相比是一个相对较低的分数。推断显示训练和推断之间的准确率差距很大可能有很多原因。一个可能的原因是全连接层没有考虑到前面截图中显示的区域信息。在深度学习中，我们使用卷积层来使网络学习空间信息。

现在，让我们使用 cuDNN 实现卷积层，将其添加到网络中，并比较模型的性能。

使用 cuDNN 的卷积神经网络

cuDNN 库为卷积操作提供了优化的性能。通过创建一个卷积层，我们将覆盖 API 的配置，用于前向和后向操作。

卷积网络层对输入数据进行卷积处理。当你想要构建一个了解区域信息的神经网络时，这种网络架构是很有用的。回想一下，在第七章中的卷积实现，CUDA 中的并行编程模式，它需要相当大的内存带宽，并需要进一步优化以获得最佳性能。然而，使用 cuDNN 库，我们也可以获得最佳性能，因为我们不必重新发明轮子。

卷积层的实现与全连接层的实现类似。然而，由于 cuDNN 库的存在，有两个不同之处：我们不必像以前那样完全实现那么多细节，我们需要为操作分配一个工作空间大小。对于每个卷积操作——前向、反向滤波器和反向输入——都需要额外的内存空间，取决于它们的算法。算法可以根据给定的输入/输出/滤波器张量维度而变化。详细的 API 调用将在稍后处理。

与其他层一样，它有三个工作阶段。对于推理阶段，我们将调用cudnnConvolutionForward()和cudnnAddTensor()。对于反向阶段，我们将调用cudnnConvolutionBackwardData()、cudnnConvolutionBackwardFilter()和cudnnConvolutionBackwardBias()。最后，对于更新阶段，我们可以重用全连接层的代码。该层的配置概述如下：

实现前向传播

在深度学习神经网络中，通常会与卷积网络一起使用池化层。池化层只是根据简单的规则选择输入数据进行输出。以下图示显示了最大池化的例子：

使用 cuDNN 库，我们将实现这两个卷积操作。

卷积层

与全连接层类似，这个卷积层有权重和偏置参数。在全连接层中，我们使用了 cuBLAS，它不需要 cuDNN 相关的描述符。然而，我们将使用 cuDNN 卷积函数，因此需要使用滤波器描述符和卷积操作描述符。以下代码显示了在构建层时应该初始化的资源：

Conv2D::Conv2D(std::string name,
        int out_channels, kernel_size, stride, padding, dilation):
        out_channels_(out_channels), kernel_size_(kernel_size),
        stride_(stride), padding_(padding), dilation_(dilation) {
  name_ = name;
  cudnnCreateFilterDescriptor(&filter_desc_);
  cudnnCreateConvolutionDescriptor(&conv_desc_);
  cudnnSetConvolution2dDescriptor(conv_desc_,
    padding_, padding_, stride_, stride_, dilation_,dilation_,
    CUDNN_CROSS_CORRELATION, CUDNN_DATA_FLOAT);
}

由于我们在模型构建时提供了卷积操作信息，我们可以指定卷积描述符。然而，滤波器的操作可以在推断时指定，因为我们可以在那时学习输入张量的大小。现在，让我们实现卷积层的前向传递。

正如我们之前讨论的，我们可以用输入张量大小初始化卷积层。这个输入张量大小会影响输出张量的大小。以下代码显示了前向传递中的参数初始化步骤：

// initialize weights and bias
if (weights_ == nullptr) {
  // initialize containers handles
  cudnnSetFilter4dDescriptor(filter_desc_, 
    CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW,
    out_channels_, input->c(), kernel_size_, kernel_size_);

  weights_ = new Blob<float>(out_channels_, input->c(), kernel_size_, kernel_size_);
  biases_ = new Blob<float>(1, out_channels_); // bias size
  bias_desc_ = biases_->tensor();
}

然后，我们需要更新输入资源，初始化输出 blob，创建 cuDNN 工作空间，并初始化权重参数，如下所示：

// initilaize input and output
if (input_ == nullptr || batch_size_ != input->n()) {
  // initialize input
  input_ = input;
  input_desc_ = input->tensor();
  batch_size_ = input->n();

  // getting output tensor size
  cudnnGetConvolution2dForwardOutputDim(
    conv_desc_, input_desc_, filter_desc_,
    &output_size_[0], &output_size_[1], 
    &output_size_[2], &output_size_[3]);

  // initialize output blob
  if (output_ == nullptr)
    output_ = new Blob<float>(output_size_);
  else
    output_->reset(output_size_);
  output_desc_ = output_->tensor();

  // initialize weights
  if (!freeze_)
    init_weight_bias();

  // initialize workspace for cudnn
  set_workspace();
}

为了获得输出张量大小，我们使用cudnnGetConvolution2dForwardOutputDim()函数。该函数根据输入张量大小、卷积操作和滤波器大小输出维度大小信息。然后，我们重用了在全连接层中使用的相同参数初始化代码。

要调用 cuDNN 的卷积 API，我们需要提供其工作算法和工作空间内存。我们这样做是因为 cuDNN 根据卷积大小选择最佳卷积算法，并且需要立即进行测量。确定算法后，cuDNN 可以确定工作空间大小。卷积层需要进行前向传播的卷积操作、输入数据的梯度和权重的梯度。我们需要分别处理每个算法，但我们可以分配一个工作空间，因为工作空间专门用于每个卷积操作。

因此，我们创建的工作空间需要具有每个卷积算法所需的最大大小。以下代码显示了我们如何使用它们并管理工作空间：

Conv2d::set_workspace() {
  size_t temp_size = 0;

  // fwd
  cudnnGetConvolutionForwardAlgorithm(cuda_->cudnn(),
    input_desc_, filter_desc_, conv_desc_, output_desc_,
    CUDNN_CONVOLUTION_FWD_PREFER_FASTEST, 0, &conv_fwd_algo_);
  cudnnGetConvolutionForwardWorkspaceSize(cuda_->cudnn(),
    input_desc_, filter_desc_, conv_desc_, output_desc_, 
    conv_fwd_algo_, &temp_size);
  workspace_size = std::max(workspace_size, temp_size);

  // bwd - data
  cudnnGetConvolutionBackwardDataAlgorithm(cuda_->cudnn(), 
    filter_desc_, output_desc_, conv_desc_, input_desc_, 
    CUDNN_CONVOLUTION_BWD_DATA_PREFER_FASTEST, 0, 
    &conv_bwd_data_algo_);
  cudnnGetConvolutionBackwardDataWorkspaceSize(cuda_->cudnn(),
    filter_desc_, output_desc_, conv_desc_, input_desc_, 
    conv_bwd_data_algo_, &temp_size);
  workspace_size = std::max(workspace_size, temp_size);

  // bwd - filter
  cudnnGetConvolutionBackwardFilterAlgorithm(cuda_->cudnn(),
    input_desc_, output_desc_, conv_desc_, filter_desc_,
    CUDNN_CONVOLUTION_BWD_FILTER_PREFER_FASTEST, 0, 
    &conv_bwd_filter_algo_);
  cudnnGetConvolutionBackwardFilterWorkspaceSize(cuda_->cudnn(),
    input_desc_, output_desc_, conv_desc_, filter_desc_, 
    conv_bwd_filter_algo_, &temp_size);
  workspace_size = std::max(workspace_size, temp_size);

  if (workspace_size > 0) {
    if (d_workspace != nullptr)
      cudaFree(d_workspace);
    cudaMalloc((void**)&d_workspace, workspace_size);
  }
}

每个卷积算法都使用单独的类型进行指定，即cudnnConvolutionFwdAlgo_t、cudnnConvolutionBwdDataAlgo_t和cudnnConvolutionBwdFilterAlgo_t。我们可以通过将它们声明为类成员变量来使用它们，即conv_fwd_algo_、conv_bwd_data_algo_和conv_bwd_filter_algo_。

现在，在初始化后，我们编写前向处理代码。我们使用滤波器进行卷积并添加偏差。以下代码显示了 cuDNN 卷积前向实现：

cudnnConvolutionForward(cuda_->cudnn(), &cuda_->one, input_desc_, input_->cuda(), \
    filter_desc_, weights_->cuda(), conv_desc_, conv_fwd_algo_, d_workspace, workspace_size, \
    &cuda_->zero, output_desc_, output_->cuda());
cudnnAddTensor(cuda_->cudnn(), &cuda_->one, bias_desc_, biases_->cuda(), \
    &cuda_->one, output_desc_, output_->cuda());

卷积的结果将使用输出 blob 传递到下一层。

实现反向传播

在反向传播中，我们应该计算偏差的梯度、权重的梯度和输入数据的梯度。为此，我们需要在第一次迭代中创建 blob 以便我们可以存储它们。它们的大小不取决于批处理大小，所以我们只需要确保它们被创建。初始化步骤可以实现如下：

// initialize grad_output back-propagation space
if (grad_weights_ == nullptr) {
  grad_output_  = grad_output;
  grad_weights_ = new Blob<float>(weights_->shape());
  grad_biases_  = new Blob<float>(1, biases_->c());
  grad_input_   = new Blob<float>(input_->shape());
}

然后，我们调用 cuDNN 反向卷积 API，如下所示：

Blob<float> *Conv2D::backward(Blob<float> *grad_output) {
  ... { initialization step } ...

  // gradients of biases
  cudnnConvolutionBackwardBias(cuda_->cudnn(),
    &cuda_->one, 
    output_desc_, grad_output->cuda(),
    &cuda_->zero, 
    bias_desc_, grad_biases_->cuda());

  // gradients of weights 
  cudnnConvolutionBackwardFilter(cuda_->cudnn(),
    &cuda_->one, 
    input_desc_, input_->cuda(), 
    output_desc_, grad_output_->cuda(),
    conv_desc_, conv_bwd_filter_algo_, d_workspace, workspace_size,
    &cuda_->zero, 
    filter_desc_, grad_weights_->cuda());

  // gradients of input data
  if (!gradient_stop_)
    cudnnConvolutionBackwardData(cuda_->cudnn(),
      &cuda_->one, 
      filter_desc_, weights_->cuda(), 
      output_desc_, grad_output->cuda(), 
      conv_desc_, conv_bwd_data_algo_, d_workspace, workspace_size,
      &cuda_->zero, 
      input_desc_, grad_input_->cuda());

然后，我们将输入数据的梯度传递给前一层以传播梯度。在更新步骤中，我们将使用基类的梯度更新代码来更新权重和偏差的梯度。在全连接层中实现反向传播时，我们已经涵盖了这一点。如果这是第一层，则我们也可以跳过计算输入数据的梯度。

使用 cuDNN 的池化层

池化层有两个特点。首先，它的输出大小与卷积层不同，cuDNN 为此提供了相应的 API。其次，它没有任何内部权重。

为了指定池化操作，我们可以使用 cuDNN 的cudnnPoolingDescriptor_t函数，并在类构造函数中创建和指定 cuDNN 的池化描述符，如下所示：

cudnnCreatePoolingDescriptor(&pool_desc_);
cudnnSetPooling2dDescriptor(pool_desc_, mode_, CUDNN_PROPAGATE_NAN,
  kernel_size_, kernel_size_, padding_, padding_, stride_, stride_);

现在，让我们实现池化层的前向和反向操作。

实现前向传播

池化层有助于减小张量的大小。因此，我们需要计算输出大小。我们可以使用cudnnGetPooling2dForwardOutputDim()函数来计算大小，就像我们在卷积层实现中所做的那样。此外，张量大小取决于批处理大小。这意味着如果批处理大小发生变化，我们需要更新张量大小。以下代码显示了我们如何初始化输入和输出 blob：

if (input_ == nullptr || batch_size_ != input->n()) {
  input_ = input;

  // resource initialize
  input_desc_ = input_->tensor();
  batch_size_ = input->n();

  // setting output
  cudnnGetPooling2dForwardOutputDim(pool_desc_, input_desc_, 
    &output_size_[0], &output_size_[1], &output_size_[2], 
    &output_size_[3]);
  if (output_ == nullptr)
    output_ = new Blob<float>(output_size_);
  else
    output_->reset(output_size_);

  output_desc_ = output_->tensor();
}

对于前向传播，我们调用cudnnPoolingForward()函数，如下所示：

Blob<float> *Pooling::forward(Blob<float> *input) {
  ... { initialization step } ...

  cudnnPoolingForward(cudnnHandle, pool_desc_, &one, 
    input_desc_, input_->cuda(),
    &zero, output_desc_, output_->cuda());
}

实现反向传播

对于反向传播步骤，我们调用cudnnPoolingBackward()函数，如下所示：

Blob<float> *Pooling::backward(Blob<float> *grad_output) {
  if (grad_input_ == nullptr)
    grad_input_ = new Blob<float>(input_->shape());

  cudnnPoolingBackward(cudnnHandle, pool_desc_,
    &one, output_desc_, output_->cuda(), 
    output_desc_, grad_output->cuda(), 
    input_desc_, input_->cuda(), 
    &zero, input_desc_, grad_input_->cuda());
}

池化层的张量形状的输入和梯度的输入是相同的，输出和梯度的输出的形状也是相同的。因此，我们可以分别重用输入和输出的张量描述符。

现在，让我们将这些集成到单个卷积层实现中。

网络配置

现在，我们将更新我们之前的网络 LeNet。网络代码可以编写如下：

Network model;
model.add_layer(new Conv2D("conv1", 20, 5));
model.add_layer(new Pooling("pool", 2, 0, 2, CUDNN_POOLING_MAX));
model.add_layer(new Conv2D("conv2", 50, 5));
model.add_layer(new Pooling("pool", 2, 0, 2, CUDNN_POOLING_MAX));
model.add_layer(new Dense("dense1", 500));
model.add_layer(new Activation("relu", CUDNN_ACTIVATION_RELU));
model.add_layer(new Dense("dense2", 10));
model.add_layer(new Softmax("softmax"));
model.cuda();

现在，我们可以开始训练和推断阶段，因为我们已经配置了我们的层，使它们彼此连接。让我们使用以下命令编译代码：

$ nvcc -run -m64 -std=c++11 -I/usr/local/cuda/samples/common/inc -gencode arch=compute_70,code=sm_70 -lcublas -lcudnn -lnvToolsExt -o train ./train.cpp ./src/layer.cu ./src/loss.cu ./src/mnist.cpp ./src/network.cpp

然后，我们可以看到训练和测试结果如下：

正如您所看到的，该网络的训练准确度和推断准确度都比仅使用全连接网络时要高。我们还可以通过查看 NVIDIA 配置文件来确认其操作，如下所示：

混合精度操作

最新的 NVIDIA GPU 支持深度学习的混合精度操作。我们不会在本书中涵盖这一点，因为它超出了我们的范围。但是，如果您希望了解更多，可以访问 NVIDIA 提供的示例，位于/usr/src/cudnn_samples_v7/conv_sample。要访问此示例，您需要从 cuDNN 网页下载示例。此示例代码显示了如何使用 cuDNN 库进行混合精度操作。

为了使 cuDNN API 与张量核心一起工作，我们需要设置数学类型，如下所示：

cudnnSetConvolutionMathType(cudnnConvDesc, CUDNN_TENSOR_OP_MATH);

然后，我们需要使用cudnnSetTensorNdDescriptor()初始化输入和输出张量的张量描述符。这为张量提供填充，以便我们获得优化的张量核心性能。

一个很好的基于 cuDNN 的实现是cudnn-training：github.com/tbennun/cudnn-training。它将 LeNet 实现为一系列 cuDNN 函数。您可以跟踪每一行，看看 CUDNN 函数是如何工作的。

如果您有兴趣使用 cuDNN 部署您的网络，请查看以下关于 GTC-CNN 推断与 cuDNN 的视频（developer.nvidia.com/gtc/2019/video/S9644/video）。这个讲座介绍了使用 cuDNN 进行 CNN 推断的有用性能优化技巧。

在深度学习训练中使用半精度需要超过 FP16 操作的利用率。我们需要在 FP16 中计算张量，同时将权重保持在 FP32 中。此外，一些操作需要 FP32。我们称之为混合精度。cuDNN 库提供了一个名为 mnistCUDNN 的混合精度推断示例。该示例显示了输入和层数据类型的转换。如果您想了解更多关于深度学习和训练中混合精度操作的信息，请阅读以下文章：devblogs.nvidia.com/video-mixed-precision-techniques-tensor-cores-deep-learning/。

现在，我们将从性能方面讨论深度学习中的其他 GPU 使用注意事项。

循环神经网络优化

RRN 允许您在深度学习中分析顺序数据。尽管该网络具有顺序依赖性，但仍有大量的优化空间。在本节中，我们将介绍其算法以及 cuDNN 如何提供优化性能。

有许多种类型的 RNN，但 cuDNN 只支持四种，即带有 ReLU 的 RNN，带有 tanh 的 RNN，LSTM 和 GRU。它们有两个输入：来自先前网络的隐藏参数和来自源的输入。根据它们的类型，它们有不同的操作。在本实验室中，我们将介绍 LSTM 操作。下图显示了 LSTM 的前向操作：

从计算的角度来看，有八个矩阵-矩阵乘法和许多逐元素操作。根据这个估计，我们可以期望 LSTM 可能是内存受限的，因为每个操作都是内存受限的。另一方面，CUDNN 提供了cudnnRNNForwardInference()和cudnnRNNFowardTraining()RNN 函数。我们将通过测量这个函数的性能和模拟 LSTM 的性能来介绍使用这个函数的好处。为了做到这一点，我们将实现一个虚拟的 LSTM 层，并将其性能与 cuDNN LSTM 函数进行比较。

为了测试目的，我们将设置超参数如下：

int mode = 2; // LSTM in CUDNN
int seq_length = 512;
int num_layers = 4;
int hidden_size = 512;
int input_size = hidden_size;
int batch_size = 32;
float dropout_rate = 0;
bool bidirectional = 0;
int persistent = 0;

序列长度或隐藏大小可能会有所不同，这取决于问题。在这个测试中，我们将使用512作为长度，在序列研究中经常使用。CUDNN API 需要更多的选项才能工作，比如 dropout 率、双向或单向以及持久 RNN。在本节中，我们只测试 vanilla LSTM。

使用 CUDNN LSTM 操作

让我们编写一些执行cudnnRNNForwardTraining()函数作为 LSTM 层的代码：

我们需要初始化输入和输出内存空间。为了执行 cuDNN 的 RNN API，我们需要使用以下变量：

// hx, cx, hy, cy, dhy, dcy, dhx, and dcs can be null.
void *x;            // input
void *hx = nullptr; // input of initial hidden state
void *cx = nullptr; // input of cell state (LSTM)

void *y;            // output
void *hy = nullptr; // output of final hidden state
void *cy = nullptr; // output of final cell state (LSTM)

void *dy;            // input of gradient 
void *dhy = nullptr; // input of final hidden state
void *dcy = nullptr; // input of final cell state (LSTM)

void *dx;            // output of gradient at the input of rnn
void *dhx = nullptr; // output of gradient at the initial hidden state
void *dcx = nullptr; // output of gradient at the initial cell state

这些变量是 LSTM 的输入和输出。为了提供输入和获取输出，我们需要分配适当的内存空间。根据 LSTM 的定义，我们需要考虑输入、输出和隐藏层的长度。这些大小可以确定如下：

int input_length = seq_length * input_size * batch_size;
int output_length = seq_length * hidden_size * batch_size;
int hidden_length = hidden_size * batch_size * num_layers;

然后，我们可以为每个项目分配内存。

现在，我们需要为 cuDNN RNN API 设置张量描述符。以下代码显示了我们应该设置的所需张量描述符：

cudnnTensorDescriptor_t x_desc[seq_length], y_desc[seq_length], \
                        dx_desc[seq_length], dy_desc[seq_length];
cudnnTensorDescriptor_t hx_desc, cx_desc;
cudnnTensorDescriptor_t dhx_desc, dcx_desc;
cudnnTensorDescriptor_t hy_desc, cy_desc;
cudnnTensorDescriptor_t dhy_desc, dcy_desc;

对于输入和输出描述符，我们需要初始化每个元素，即批量大小和其输入大小。其他隐藏的张量描述符是用层数、批量大小和隐藏大小进行初始化的。本节不涵盖如何编写初始化代码。但是，如果您想了解更多信息，可以查看10_deep_learning/03_rnn文件中的代码。

我们还需要为 RNN 操作提供一个工作空间，就像我们为卷积操作做的那样：

void *workspace;
cudnnFilterDescriptor_t w_desc, dw_desc;
cudnnSetRNNDescriptor_v6(cudnnHandle, rnn_desc,
                         hidden_size, num_layers, dropout_desc, CUDNN_LINEAR_INPUT,
                         bidirectional ? CUDNN_BIDIRECTIONAL : CUDNN_UNIDIRECTIONAL,
                         CUDNN_LSTM, CUDNN_RNN_ALGO_STANDARD, CUDNN_DATA_FLOAT));
size_t weight_size;
cudnnGetRNNParamsSize(cudnnHandle, rnn_desc, x_desc[0], &weight_size, CUDNN_DATA_FLOAT);
cudaMalloc((void**)&workspace, weight_size);

然后，我们可以根据工作空间的大小设置滤波器描述符，如下所示：

dimW = {weight_size / sizeof(float), 1, 1}
cudnnCreateFilterDescriptor(&w_desc);
cudnnCreateFilterDescriptor(&dw_desc);
cudnnSetFilterNdDescriptor(w_desc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, 3, dimW);
cudnnSetFilterNdDescriptor(dw_desc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, 3, dimW);
cudnnRNNForwardTraining(cudnnHandle, rnn_desc, seq_length,
                x_desc, x, hx_desc, hx, cx_desc, cx,
                w_desc, w, 
                y_desc, y, hy_desc, hy, cy_desc, cy,
                workspace, workspace_size, reserved_space, 
                reserved_size);

我们可以使用cudaEvnetRecoard()和 flops 计算来衡量它们的性能。例如，前向操作可以配置为以下方程：

然后，我们将通过将批量大小从 32 增加到 256 来测试我们的实现，每次增加 32。适用的测试范围可能会有所不同，以及 GPU 的内存大小。

在本节中，我们实现了基于 LSTM 的模拟和cudnnRNNForwardTraining()调用。我们部分模拟的版本只有 GEMM 操作，这是最计算密集的操作。现在，让我们比较这些实现的性能。

实现虚拟 LSTM 操作

在我们的实现中，我们将专注于模拟 LSTM 的主要操作，而不是完全实现它。

让我们确定 LSTM 网络的超参数。一般来说，输入序列长度范围从 512 到 2,048。层数的数量是不确定的。但是，由于tanh操作，它不能太大。对于输入大小，我们将使用 512。通常情况下，批量大小在 RNN 使用方面在 32 到 256 之间。CUDNN 需要更多关于 dropout 率、双向或单向以及是否使用持久 RNN 的输入。我们现在不使用它们。我们的 LSTM 配置信息如下：

现在，我们将部分实现 LSTM 操作以测量计算强度。正如我们之前讨论的，LSTM 有两个矩阵-矩阵乘法需要计算。LSTM 操作将为输入序列的每个元素以及每个层计算。然后，操作可以配置如下：

for (int layer = 0; layer < num_layers; layer++) {
  for (int linear_layer = 0; linear_layer < 4; linear_layer++) {
    for (int sequence = 0; sequence < seq_length; sequence++) {
      cublasSgemm(cublas_handle, CUBLAS_OP_T, CUBLAS_OP_N,
      hidden_size, input_size, batch_size,
      &alpha, input_weight, input_size, x, input_size,
      &beta, h, hidden_size);
      cublasSgemm(cublas_handle, CUBLAS_OP_T, CUBLAS_OP_N,
      hidden_size, hidden_size, batch_size,
      &alpha, recurrent_weight, hidden_size,
      h, hidden_size,
      &beta, y, hidden_size);
    }
  }
}

我们可以使用更多的逐元素操作，但这只是近似计算强度，所以我们暂时不考虑它们。

比较 CUDNN 和 SGEMM LSTM 的性能

让我们比较它们的性能以及不同的批处理大小，如下所示的代码实现在main()函数中：

for (int step = 1; step <= 8; step++)
{
 batch_size = 32 * step;
 printf("Batch Size: %3d\n", batch_size);
 rnn_operation(seq_length, num_layers, hidden_size, input_size,   
   batch_size, dropout_rate, bidirectional, mode, persistent);
 cublas_operation(mode, 2ull, input_size, hidden_size, seq_length, batch_size, num_layers);
}

然后，我们可以使用以下命令编译和执行示例源代码：

$ nvcc -run -m64 -std=c++11 -I/usr/local/cuda/samples/common/inc -gencode arch=compute_70,code=sm_70 -lcublas -lcudnn -lcurand -o rnn ./rnn.cpp

以下图表显示了来自 Tesla V100 卡的 cuBLAS 和 cuDNN 的性能：

在上图中，两种实现在性能上有很大差异。cuDNN 的 LSTM 性能比使用 cuBLAS 模拟的 LSTM 要好得多。此外，LSTM 操作的性能遵循 Tesla V100 GPU 的屋顶线。另一方面，两个 SGEMM 操作并没有显示出这种性能，因为矩阵大小不够大以获得完整的性能。要从 Tesla V100 获得 10 TFlops，矩阵大小应与 1,024 的平方相似或更大。然而，正如我们所看到的，我们的矩阵大小大约是 512 的平方。

LSTM 优化在以下 NVIDIA 文章中有解释：devblogs.nvidia.com/optimizing-recurrent-neural-networks-cudnn-5。它结合了矩阵-矩阵乘法，融合逐元素操作，多个流和多层并行化。

RNN 的优化版本之一是持久 RNN（svail.github.io/persistent_rnns），由 Greg Diamos 介绍。尽管他的实现不包括 LSTM 和 GRU，但您可以了解 RNN 如何进行优化。

深度学习框架的性能分析

一般来说，我们使用 TensorFlow、PyTorch 和 MxNet 等深度学习框架开发和研究神经网络。由于这些框架，我们可以有效地开发复杂的模型。然而，当涉及性能工程时，由于性能分析工具的能力，理解框架下 GPU 操作是一个陡峭的学习曲线。例如，使用 Chrome 跟踪进行性能分析在模型简单时很有用，但在模型复杂时就不那么有用。

在第五章中，CUDA 应用程序性能分析和调试，我们介绍了NVIDIA 工具扩展（NVTX），它允许我们在 GPU 应用程序中进行自定义注释，并使用 NVIDIA Nsight Systems 查看时间轴。对于复杂的应用程序，程序员分析其性能并找到瓶颈非常有用。

在本节中，我们将介绍如何通过修改 ResNet-50 示例代码在 PyTorch 和 TensorFlow 中使用 NVTX。示例代码可以在本书的 GitHub 存储库的10_deep_learining/05_framework_profile文件夹中找到。您可以从github.com/nvidia/DeepLearningExamples获取原始源代码。

为了简化工作环境配置，我们将使用NVIDIA GPU 云（NGC）深度学习容器用于 PyTorch 和 TensorFlow。如果您需要了解 NGC 或容器的基本用法，请访问本书附录中的 NGC。

现在，让我们先从 PyTorch 开始。

对 PyTorch 模型进行性能分析

在 PyTorch 中，我们可以使用torch.cuda.nvtx.range_push("foo")和torch.cuda.nvtx.range_pop()来放置自定义标签。这保持了原始的 CUDA NVTX API，即nvtxRangePush()和nvtxRangePop()。让我们看看 NVTX 注释如何帮助我们在时间轴上理解深度学习操作。在接下来的步骤中，我们将使用05_framework_profile/pytorch/RN50v1.5文件中的 ResNet-50 示例代码：

我们将在train()函数中的训练循环中放置 NVTX 注释以注释step值。该函数可以在image_classificaiton/training.py文件中找到。以下截图显示了训练循环和分别在第 234 行和第 260 行的 NVTX 注释：

在上述代码中，训练操作是在step函数中实现的，该函数由get_train_step()函数定义。因此，我们需要在该函数中放置 NVTX 注释以了解更多信息。

让我们在第 164 行的get_train_step()函数中添加一些 NVTX 注释。该函数返回_step()函数，其中包括训练操作。因此，我们将在该函数中放置 NVTX 注释。训练过程包括前向和反向传播、全局归约和优化（更新权重）。以下截图显示了在第 166 行和第 171 行的前向传播的注释：

通过这种方式，我们可以在其余操作上放置其他注释。

我们还可以为模型层添加 NVTX 注释。在这个例子中，ResNet-50 模型是在image_classification/resnet.py文件中实现的。以下截图显示了网络的示例注释：

正如我们所看到的，我们可以按照 ResNet 架构放置 NVTX 注释。如果我们在每个构建块中放置注释，我们可以获得更多信息。

现在，让我们对模型进行分析。正如我们之前讨论的，我们将使用 NGC 深度学习容器，即 PyTorch。imagenet数据集位于/raid/datasets/imagenet/raw-data文件夹中。为了限制分析时间范围，我们将使用延迟选项（-y）和持续时间选项（-d）。以下代码显示了一个执行容器并对网络进行分析的 bash shell 脚本：

#/bin/bash

CODE_PATH="RN50v1.5"
DATASET_PATH="/raid/datasets/imagenet/raw-data/"
OUTPUT_NAME="resnet50_pyt"

# default profile
docker run --rm -ti --runtime=nvidia \
    -v $(pwd)/${CODE_PATH}:/workspace \
    -v ${DATASET_PATH}:/imagenet \
    nvcr.io/nvidia/pytorch:19.08-py3 \
       nsys profile -t cuda,nvtx,cudnn,cublas -o ${OUTPUT_NAME} 
         -f true -w true -y 60 -d 20 \
       python /workspace/main.py --arch resnet50 -b 64 
         --fp16 /imagenet

执行后，上述代码将在 RN50v1.5 目录中生成 profiled 结果，即resnet50_pyt.qdrep。

最后，使用 NVIDIA Nsight Systems 打开 profiled 输出resnet50_pyt.qdrep，并查看操作。以下截图显示了带有 NVTX 注释的测量步骤：

在这里，我们可以看到反向操作所花费的时间是前向操作的两倍。此外，PyTorch 将主机线程分开用于训练循环和反向传播。从内核分析来看，耗时最长的点是逐元素的内核执行。让我们扩大前向传递以查看层的执行时间，如下截图所示：

在这里，我们可以看到第二个卷积块需要最长的时间来完成。如果这一层存在效率低下的点，我们可以进一步挖掘。如果某个操作被确定为瓶颈并需要优化，我们还可以使用 NVIDIA Nsight Compute 来分析特定的内核函数。比较主机 API 跟踪和 GPU，我们可以看到时间持续时间是不同的。这是因为主机和 GPU 操作是异步的。因此，当我们从主机测量 GPU 执行时间时，我们需要谨慎。现在，让我们看一下优化步骤，如下截图所示：

我们可以看到，从主机和 GPU 的测量执行时间中存在巨大差异。主机的测量执行时间为 25.367 毫秒，而 GPU 的时间为 4.048 毫秒。其操作主要是逐元素操作，其执行被延迟直到反向传播完成。我们还可以找到异步执行。之后，我们可以看到cudaDeviceSynchronize()函数，该函数防止当前步骤被下一步骤更新。

我们还可以通过设置环境来禁用这些异步操作，即CUDA_LAUNCH_BLOCKING=1。我们可以使用环境选项（-e）将其传递给 Nsight System 的配置选项。然后，我们可以分析应用程序的align操作与主机和内核函数。

PyTorch 在其 CUDA 对象中具有几个具有 NVTX 特色的 API。 PyTorch 文档可以在pytorch.org/docs/stable/_modules/torch/cuda/nvtx.html找到。通过直接在 PyTorch 中调用 NVTX API，将调用 CUDA NVTX API。这意味着我们可以在分析时间线中获得自定义标记的 NVTX 标记。

对 TensorFlow 模型进行分析

对 TensorFlow 图进行分析需要使用启用 NVTX 注释的 NVTX 插件。要在 TensorFlow 中使用 NVTX 注释，我们需要使用以下命令安装nvtx-plugins-tf Python 插件：

$ pip install nvtx-plugins-tf

但是，如果我们使用的是版本晚于 19.08 的 NGC TensorFlow 容器，则无需执行此操作。

TensorFlow 图形 API 是符号 API，因此它们需要特定的编程方法。 NVTX 插件为此提供了两个选项：装饰器和 Python 函数。

以下是 NVTX 装饰器的示例：

import nvtx.plugins.tf as nvtx_tf
ENABLE_NVTX=true
@nvtx_tf.ops.trace(message='Dense Block', domain_name='Forward',
        grad_domain_name='Gradient', enabled=ENABLE_NVTX, 
        trainable=True)
def dense_layer(x):
    x = tf.layers.dense(x, 1000, activation=tf.nn.relu, name='dense_1')
    x = tf.layers.dense(x, 1000, activation=tf.nn.relu, name='dense_2’) 
return x

以下是 NVTX Python 函数的示例：

import nvtx.plugins.tf as nvtx_tf
ENABLE_NVTX=true
x, nvtx_context = nvtx_tf.ops.start(x, message='Dense Block', \ 
        domain_name='Forward’, grad_domain_name='Gradient’, 
        enabled=ENABLE_NVTX, trainable=True)
x = tf.layers.dense(x, 1000, activation=tf.nn.relu, name='dense_1')
x = tf.layers.dense(x, 1000, activation=tf.nn.relu, name='dense_2’) 
x = nvtx_tf.ops.end(x, nvtx_context)

NVTX 插件提供了 NVTXHook，它允许我们对 TF 估算器和会话进行分析。例如，我们可以按以下方式使用该钩子：

from nvtx.plugins.tf.estimator import NVTXHook

nvtx_callback = NVTXHook(skip_n_steps=1, name='Train’)
training_hooks=[]
training_hooks.append(nvtx_callback)

然后，我们可以使用以下代码将其应用于任一选项：

with tf.train.MonitoredSession(hooks=training_hooks) as sess:

或者，我们可以使用以下代码：

tf.estimator.Estimator(hooks=training_hooks, ...)

现在，让我们将其应用到示例 ResNet-50 代码中并进行操作审查。示例代码可以在05_framework_profile/tensorflow/RN50v1.5文件夹中找到：

让我们首先将NVTXHook应用于估算器。训练图的定义可以在runtime/runner.py文件的第 312 行找到。在构建图之前，我们将NVTXHook附加到钩子列表中，如下面的代码块所示：

然后，我们将 NVTX 注释应用于模型构建函数。model_build()函数可以在model/resnet_v1_5.py文件的ResnetModel类中找到。以下代码显示了如何在model_build()函数中的conv1层上使用 Python 函数放置 NVTX 注释的示例：

在上述代码中，当使用nvtx_tf.ops.start()和nvtx_tf.ops.end()函数时，我们需要谨慎选择适当的输入和输出。只在其他层中放置 NVTX 注释。确保最终的全连接层输出是网络的输出。

我们还必须禁用用于检查可训练变量数量的代码。如果 NVTX 的trainable参数值为True，则大小会发生变化。在resnet_v1_5.py文件的第 174 行，有一段断言代码，用于检查该变量的数量。只需将其注释掉，如下所示：

我们还使用 NVTX 装饰器来构建 ResNet 模块。在model/blocks目录中，我们可以在conv2d_blocks.py和resnet_bottleneck_block.py中找到conv2d和 ResNet 瓶颈块的实现。在conv2d_blocks.py文件中，我们可以装饰conv2d_block()函数以注释 NVTX 分析，如下所示：

同样，我们也可以对resnet_bottleneck_block.py文件执行相同操作：

现在，让我们对模型进行性能分析。就像我们使用 PyTorch 容器一样，我们将使用 TensorFlow 的 NGC 容器。我们假设imagenet数据集的tfrecord文件位于/raid/datasets/imagenet/tfrecord目录中。以下代码显示了一个执行容器并对网络进行性能分析的 bash shell 脚本：

#/bin/bash

CODE_PATH="RN50v1.5"
DATASET_PATH="/raid/datasets/imagenet/tfrecord"
OUTPUT_NAME="resnet50_tf"

# default profile
docker run --rm -ti --runtime=nvidia \
    -v $(pwd):/result \
    -v $(pwd)/${CODE_PATH}:/workspace \
    -v ${DATASET_PATH}:/imagenet \
    nvcr.io/nvidia/tensorflow:19.08-py3 \
        nsys profile -t cuda,nvtx,cudnn,cublas -o ${OUTPUT_NAME} 
                     -f true -w true -y 40 -d 20 \
            python /workspace/main.py --mode=training_benchmark 
                                      --warmup_steps 200 \
                --num_iter 500 --iter_unit batch 
                --results_dir=results --batch_size 64

当我们执行这个函数时，我们将在RN50v1.5目录中得到resnet50_tf.qdrep文件。

最后，让我们使用 NVIDIA Nsight System 审查分析输出：

在这里，我们可以确认反向传播所花费的时间是前向传播的两倍。这个示例代码与 CPU 和 GPU 不同步。因此，我们可以看到主机和 GPU 之间的时间差异更大。当我们在构建块中放置额外的注释时，我们将能够在层中看到子块的注释。

使用 NVIDIA Nsight Systems 进行性能分析在多 GPU 训练中监视所有归约操作的执行时间时提供了额外的好处。以下截图显示了一个使用两个 GPU 进行训练的 GPU 的性能分析结果：

在突出显示的行中，我们可以看到ncclAllRecude()函数，它同时调用了反向传播。通过这样做，我们不会延迟所有归约操作。这个示例代码使用 Horovod 来训练多个 GPU。如果你想了解更多，请访问 Horovod 的 GitHub 页面：github.com/horovod/horovod。你可以从这里获取文档和示例代码。

总结

在本章中，我们学习了如何使用 CUDA 库进行深度学习和性能优势。在回顾它们的用途时，我们将它们与每个步骤的深度学习机制进行匹配。由于我们可以使用的深度学习库，我们可以实现一个简单的 CNN，而不必实现算法。然后，我们使用 NVTX 注释在 PyTorch 和 TensorFlow 中对 ResNet-50 模型进行了性能分析。

对于一些深度学习工程师和研究人员来说，实现基本算法可能是不切实际的。然而，了解性能因素和基本操作可以帮助您构建高效和有效的基于深度学习的产品。如今，我们看到许多产品化的基于深度学习的服务。工程师们花费大量资源将他们训练好的模型产品化，以及训练他们的模型，以便获得尽可能低的错误率。希望您能够了解如何在深度学习应用中使用 NVTX 性能分析。利用这些知识，您可以更好地利用您的 GPU。祝你好运！

附录

CUDA 是一个并行编程平台。学习 CUDA 不仅意味着学习语言，还意味着具有一些与 GPU 相关的工程技能。这个工程领域可以是监控、环境设置、性能理解、容器化等等。本章提供了一些提示，以帮助工程师使用 GPU。我们可以涵盖更多的主题，但以下主题对于那些想要学习 CUDA 及其 GPU 操作的人来说将是有帮助的。

在本章中，我们将涵盖以下主题：

有用的nvidia-smi命令
Windows 中的 WDDM/TCC 模式
性能建模
探索基于容器的开发

有用的 nvidia-smi 命令

在本节中，我们将涵盖nvidia-smi的监控功能和管理操作。nvidia-smi是NVIDIA 管理库（NVML）的命令行接口（CLI）。该库使得对 NVIDIA 设备进行管理和监控成为可能。nvidia-smi还通过该库提供了对设备的直接查询和命令。数据以纯文本或 XML 格式通过stdout或文件呈现。它提供了几个管理工具，用于更改设备统计信息。

nvidia-smi是一个包装 NVML C/C++ API 的 CLI 应用程序。它通过 NVML 从 NVIDIA 驱动程序获取请求的信息。NVML 还提供了用于其他语言（如 Python 和 Perl）的 API。

基本上，nvidia-smi为用户报告了以下已安装的 GPU 统计信息：

第一行报告了驱动程序版本和支持的 CUDA 版本
第二行显示了 GPU 的统计格式
每个连续的行包含每个 GPU 的统计信息，包括以下内容：
GPU ID
操作模式：
持久模式（开/关）
Tesla 计算集群（TCC）/Windows 显示驱动模型（WDDM）模式
风扇速度
GPU 温度
性能模式
功率使用和容量
总线 ID
内存使用和已安装内存
计数的纠错码（ECC）
GPU 利用率
计算模式

基本上，nvidia-smi可以处理包括 Tesla、Quadro 和 GeForce 在内的所有 NVIDIA GPU 卡。启用的功能可能因型号和类型而异。例如，ECC 错误计数在 Tesla 和 Quadro 卡中可用，而在 GeForce 中不可用，因为它在设备内存中不提供 ECC 功能。

nvidia-smi报告的格式在各个操作系统上都是相同的。以下截图显示了 Windows 的输出：

以下截图显示了 Linux 的输出：

因此，我们可以阅读报告并以相同的格式设置 GPU 操作。现在，让我们继续看一下经常使用的命令。默认的nvidia-smi CLI 的用法如下：

$ nvidia-smi [option1 [arg]] [option2 [arg]] ...

首先，根据监视目的，经常使用以下选项：

-i，--id=：用于选择目标 GPU
-l，--loop=：以指定的秒间隔报告 GPU 的状态
-f，--filename=：用于记录到指定文件中

此列表涵盖了可以帮助我们从 GPU 获取详细信息的nvidia-smi选项。

获取 GPU 的信息

当我们使用--query（-q）选项时，nvidia-smi报告结构化输出。因此，我们可以了解收集了哪些信息。我们可以获取 GPU 的利用率、功率、内存和时钟速度统计信息。另一方面，如果我们希望连续监视 GPU 的状态，这种格式就不太有用了。

获取格式化信息

我们需要监视的基本 GPU 统计信息是功率、温度、核心利用率和内存使用情况。这可以很容易地通过--query-gpu命令完成：

$ nvidia-smi --query-gpu=timestamp,name,pci.bus_id,driver_version,pstate,pcie.link.gen.max,pcie.link.gen.current,temperature.gpu,utilization.gpu,utilization.memory,memory.used,memory.free,memory.used --format=csv -l 1

以下命令显示了我们可以使用的一些选项，以检测时钟调节的性能降低原因：

$ nvidia-smi --query-gpu=index,clocks_throttle_reasons.active,clocks_throttle_reasons.gpu_idle,clocks_throttle_reasons.applications_clocks_setting,clocks_throttle_reasons.sw_power_cap,clocks_throttle_reasons.hw_slowdown,clocks_throttle_reasons.hw_thermal_slowdown,clocks_throttle_reasons.hw_power_brake_slowdown,clocks_throttle_reasons.sync_boost --format=csv

GPU 时钟调节的原因可能是功率限制、过热和同步提升。功率限制意味着 GPU 的功耗受用户设置或系统中电源供应商的性能限制。过热也是由于散热环境不佳而频繁出现的调节原因。

功耗管理模式设置

您可以使用以下命令找出每个 GPU 的最大功耗：

$ nvidia-smi -i <device id> -pl N

设置 GPU 的时钟速度

默认情况下，GPU 的时钟速度会根据需求变化，以节省功耗，最大化功耗效率。为了最大化 GPU 的性能并减少延迟，特别是在基准测试情况下，我们可以确保 GPU 具有最大时钟速度并禁用 GPU 驱动程序。

首先，我们需要将 GPU 设置为持久模式。这样做意味着 GPU 驱动程序模块始终加载到内核中，并减少了初始响应时间。这个选项只在 Linux 上可用，因为 Windows 不会卸载 GPU 驱动程序。持久模式设置命令如下：

$ sudo nvidia-persistenced

然后，我们可以设置最大支持的时钟。这个值会根据您使用的 GPU 而有所不同：

$ nvidia-smi -q -d SUPPORTED_CLOCKS
$ sudo nvidia-smi -ac <Mem clock, Graphics clock>

例如，Tesla V100 卡可以使用以下命令进行设置：

$ sudo nvidia-smi -ac 877,1380 # V100 PCIe
$ sudo nvidia-smi -ac 877,1530  # V100 SMX

GPU 设备监控

这个命令每秒探测一次所选 GPU 的设备状态：

$ nvidia-smi dmon -s pucvmet -i -0

以下截图显示了前面命令的结果。我们正在监控的设备状态显示为0：

收集到的信息可以使用-s选项指定，如下所示：

p: 功耗和温度
u: 利用率
c: 处理器和内存时钟
v: 功耗和温度违规
m: FB 和 Bar1 内存
e: ECC 错误和 PCIe 重播错误
t: PCIe 接收和发送吞吐量

监控 GPU 利用率以及多个进程

如果您在单个 GPU 上使用多个进程操作，可以考虑使用这个命令。这个命令收集 GPU 统计信息，以及它们正在使用的进程。这意味着您可以确定哪个进程被 GPU 共享限制，内存时序的空间等等：

$ nvidia-smi pmon -i 0 -s u -o T

以下截图显示了带有进程 ID（PID）的nvidia-smi输出，这有助于确定哪个进程正在使用哪个 GPU 资源：

前面截图中的每一列显示了每个 GPU 的计算单元利用率或内存使用情况：

sm%: CUDA 核心利用率
mem%: 内存操作的采样时间比率
enc%/dec%: 硬件编码器利用率
fb: FB 内存使用

获取 GPU 拓扑信息

在多 GPU 系统中，使用nvidia-smi获取 GPU 拓扑信息非常有用。以下命令是一个显示多 GPU 系统 GPU 拓扑的nvidia-smi命令：

nvidia-smi topo -m

以下截图显示了nvidia-smi的输出，显示了系统的拓扑结构。DGX Station 的结果是我们有四个支持 NVLink 的 V100 GPU：

根据这个结果，我们可以确认系统的 GPU 拓扑如下：

以下命令识别了 GPU 之间的点对点可访问性。我们在第六章 可伸缩多 GPU 编程中使用了这个命令：

$ nvidia-smi topo -p2p rwnap

以下是一个拥有四个 GPU 的系统的nvidia-smi拓扑输出：

点对点访问是可伸缩性或操作的重要因素。这个命令可以帮助您确认 GPU 和您的系统是否支持 GPU 之间的点对点访问。

在 Windows 中的 WDDM/TCC 模式

在 Windows 平台上，NVIDIA GPU 有两种模式：WDDM 和 TCC。WDDM 是视频卡的图形驱动程序，因此它可以渲染桌面和应用程序。如果安装的 GPU 仅用于计算，则显示渲染是无用的开销。在这种情况下，NVIDIA GPU 可以切换到仅专注于计算的模式。这种模式称为 TCC 模式。

WDDM 允许 NVIDIA GPU 与 Windows 的 WDDM 驱动程序合作，用于显示。支持 WDDM 模式是 Windows 图形的要求。另一方面，TCC 模式只用于计算。根据您的 GPU 产品和配置，GPU 的模式可以更改。

操作模式遵循四个 NVIDIA 产品类别，并且其默认模式可以变化，如下所示：

GeForce：仅支持 WDDM 模式。
Quadro/Titan：默认情况下为 WDDM 模式，但也可以在 TCC 模式下使用。
Tesla：通常默认为 TCC 模式。
Tegra：仅支持 Linux。没有 WDDM/TCC 问题。

WDDM 模式支持 CUDA 操作和使用 Nsight 调试 CUDA 应用程序，同时也支持显示。作为单个主机机器，您可以做 GPU 能做的一切。但是，TCC 模式禁用了图形驱动程序上的图形，并将 GPU 作为计算加速器启用。换句话说，当显卡不必提供显示时应使用此模式。

TCC 模式在 CUDA 处理中比 WDDM 模式具有一些优势，如下所示：

用于大规模计算
忽略 Windows 的显示超时间隔（通常为两秒），以启用长于两秒的内核操作
减少 Windows 上 CUDA 的核心启动开销
支持在 Windows 远程桌面服务中进行 CUDA 处理
使非 NVIDIA 集成图形的 NVIDIA GPU 可用，以便您可以保存全局内存

因此，如果 GPU 不提供显示服务，TCC 模式为 GPU 作为加速器带来了最佳配置。

设置 TCC/WDDM 模式

要更改 TCC 或 WDDM 模式，请使用nvidia-smi实用程序，如下所示：

$ sudo nvidia-smi -dm {0|1}

0表示 WDDM 模式，1表示 TCC 模式。

如果要为所选的 GPU 设置 TCC 模式，请使用-g选项指定目标 GPU：

$ nvidia-smi -g {GPU_ID} -dm {0|1}

当您想要将 GPU 用途分开为显示和计算时，此选项非常有用。应用这些设置后，您可能需要重新启动您的机器以应用这些更改。

我们可以通过使用nvidia-smi来确定 TCC 模式是否已启用。以下截图显示了 TCC 中的 GPU 操作模式：

通过查看第一列中 GPU 名称的右侧，我们可以确认 TCC 模式已启用。

性能建模

了解应用程序/算法和 GPU 硬件的特性以设定实际加速目标非常重要。这可以通过增加并行性来实现。我们还需要确定在优化应用程序时是否有优化 GPU 的空间。

一个简单的方法是应用阿姆达尔定律。我们可以预测应用程序中可实现的性能增益受到代码顺序部分的限制。例如，只有 50%的代码可以并行执行，而其余部分是顺序的（例如从文件中读取）。如果是这种情况，那么可以实现的最大加速比为 2 倍；也就是说，程序只能运行两倍快。然而，这种性能建模只显示了最大加速比。我们不得不假设我们可以完全并行化并消除代码并行部分的执行时间。

另一种性能建模实践是基于目标架构的性能限制因素进行分析。实际上，我们有硬件规格，其操作引入了不可避免的性能限制。通过分析这些限制，我们可以确定是否有优化的空间，并查看下一组优化策略。

Roofline 模型

每个核函数可以归类为以下类别之一：

计算受限：核心函数对每个读取或写入的数据字节进行更多的算术运算。这些应用程序需要硬件提供更多的计算 FLOPS。
内存受限：应用程序大部分时间都在读写内存，而不是计算。应用程序受到系统内存带宽的影响最大，而不是硬件的 FLOP 评级。
延迟受限：核心函数的 CUDA 线程大部分时间都在等待而不是执行。出现这种情况有很多原因。主要原因是并行性水平不佳或内存和计算资源的使用不佳。

由于所有这些限制都是由硬件引入的，我们可以绘制目标硬件的峰值性能和内存带宽以及它们的算术强度。性能曲线受硬件的峰值性能限制。我们在第三章中简要提到了这一点，CUDA 线程编程，以确定下一个优化策略。以下插图在第三章中使用，CUDA 线程编程，展示了屋顶线模型的一个示例：

为了进行任何计算，数据需要从内存传输到算术单元。它遍历不同级别的内存层次结构，峰值内存带宽取决于内存类型。算法的峰值性能可以根据其算术强度进行分类。这种强度由计算数据与加载数据的量确定。此外，这些延迟限制因素引入了计算上限。通过针对硬件规格的性能和分析，我们可以确认目标算法是否达到了峰值性能或受到了内存或延迟的限制。在任何情况下，我们都可以确定下一步。在第三章中，CUDA 线程编程，我们深入探讨了这一点。在本节中，我们将通过一个示例来关注屋顶线模型，并看看它有多有用。

屋顶线模型考虑了应用程序的操作强度。简单来说，这意味着每个操作都是从主存储器（DRAM）中的字节进行的。虽然还有更复杂的模型考虑了缓存到处理器的传输，但屋顶线模型更专注于从 DRAM 到缓存的数据传输，因此更专注于 CUDA 核心在特定 GPU 架构上所需的 DRAM 带宽。

屋顶线模型陈述如下：

"可达性能（GFLOP/s）= min（峰值浮点性能，峰值内存带宽*操作强度）"

分析雅可比方法

让我们尝试理解这个公式，并得到 V100 GPU 卡的屋顶线模型。V100 GPU 的规格如下：

80 个 SM，每个 SM 有 32 个 FP64 核心
900GB/s 的总带宽
L2 缓存：6MB
L1 缓存：10MB
寄存器：每个 SM 62KB

让我们尝试分析一个简单的雅可比方法：

for (int iy = 1; iy < NoRows; iy++)
{
    for ( int ix = 1; ix < NoCols; ix++)
    {
        Anew[ix][iy] = rhs[iy∗nx+ix] 
                     - 0.25f*(Aref[ix-1][iy] + Aref[ix+1][iy] 
                     + Aref[ix][iy-1] + Aref[ix][iy+1]);
    }
}

让我们分析上述代码的数据传输：

向量（Anew，rhs，Aref）的内存加载：I[Load] = NoRow * NoCol * 3 * 8 字节（双精度）
向量（Anew）的存储：I[store] = NoRow * NoCol * 8 字节
浮点运算：I[FP] = NoRow * NoCol * 6 FLOP

以下图表显示了 Tesla V100 卡的屋顶线分析和雅可比方法的算术强度：

V100 上雅可比的算术强度将是I[FP]/(I[Load]+I[Strore]) = 0.18 FLOP/字节

Roofline 模型清楚地显示了算法是内存绑定的，最大可达性能仅为 0.18 FLOP/字节，因此将无法达到 V100 的峰值 FLOP 评级，即 7.8 TFLOPS。然而，我们也可以通过重用获取的数据来预测优化后的可达性能。

Roofline 模型有助于根据硬件特性定义算法的上限性能。

雅各比方法

这是一种用于解决线性方程组的迭代算法。其基本操作和 GPU 优化在www.olcf.ornl.gov/wp-content/uploads/2016/01/Introduction-to-Accelerated-Computing-with-OpenACC-Jeff-Larkin.pdf中有解释。

探索基于容器的开发

开发人员和维护集群的 IT 管理员面临的一个关键挑战是软件堆栈的复杂性。每个应用程序/框架都有许多依赖关系。当这些依赖关系是不同版本时，复杂性会增加。例如，在 DL 中，Caffe 对 cuDNN 和 Python 的版本有不同的要求，与 TensorFlow 不同。在特定的组织/学院中，有许多用户，每个用户可能使用相同框架的不同版本。安装所有正确的依赖关系并设置正确的环境会导致生产力的损失。花在安装上的时间比实际工作的时间更多。另一个面临的挑战是，即使在相同的系统上运行，由于依赖关系不匹配，不同的个体几乎不可能复制结果/性能数字。例如，GROMACS 分子动力学框架有许多设置，比如编译多线程或消息传递接口（MPI）支持，MPI 的版本和 MPI 类型。特别是在人工智能领域，另一个挑战是，你能想到的每个软件框架都在快速发展，并且经常添加新的补丁。

容器为这些问题提供了解决方案。使用容器的主要优势如下：

隔离：容器为应用程序提供环境隔离
随处运行：容器提供了在不同环境中共享和测试应用程序的简单方法
轻量级：与基于虚拟机的解决方案相比，容器轻量级，并且几乎没有延迟和开销

两个最著名的容器环境是 Docker 和 Singularity。它们都有各自的优缺点。但请注意，本节不是 Docker 或 Singularity 的详尽指南。

开发人员通常会创建容器并将其发布到网上供他人使用。我们将详细解释由 NVIDIA 维护的名为Nvidia GPU Cloud（NGC）的一个这样的存储库。NGC 就像是一个为流行的深度学习（DL）、高性能计算（HPC）和虚拟现实（VR）框架提供容器的存储库。NVIDIA 会在不同的 GPU 环境中测试这些应用程序，并在向公众提供之前经过广泛的质量保证过程。这意味着性能是有保证的。

NGC 的类比是 Android 应用商店，为可以在运行 Android 操作系统的不同手机上运行的不同应用程序提供了存储库。这些应用程序经过验证并经过质量保证流程。NGC 的名称有时会让人们感到困惑，开发人员认为它是一个云。应明确指出，它是一个容器的存储库，可以被拉取到具有 GPU 的系统中并在本地运行。该容器可以在具有 GPU 的不同系统上运行，就像它可以在具有 NVIDIA Titan 卡的台式机、具有 Tesla V100 卡的服务器或 NVIDIA AI 超级计算机 DGX 上运行一样。NGC 容器也可以在 AWS 和 Azure 等云平台上运行。

主机机器的 NGC 配置

以下步骤涵盖了如何配置 NGC 工作环境以及在 NGC 中查找可用的映像：

基本安装：要在 GPU 系统上使用容器，您需要安装以下内容：

- Nvidia 驱动程序
Docker
nvidia-docker

nvidia-docker是一个开源项目，它将 NVIDIA 组件和模块加载到容器中。它基本上是 Docker 的包装器。您可以在github.com/nvidia/nvidia-docker/wiki/Installation-(version-2.0)上下载并查看安装说明。

访问 NGC 网站：现在，您可以转到 NGC 网站选择一个容器（nvidia.com/ngc），如下面的屏幕截图所示：

如您所见，容器有六个类别。选择与您相关的类别。NGC 的早期版本要求用户注册，但最近取消了此要求。

NGC 容器的基本用法

在本节中，我们将介绍如何从 NGC 注册表中拉取容器以及如何自定义我们自己的容器。这与使用 Docker 没有区别，只是我们可以访问 NGC 注册表nvcr.io。如果您已经熟悉 Docker 命令，可以跳过本节。

以下步骤解释了如何在本地 Linux 机器上的终端会话中获取和启动 NGC 容器：

找到您想要使用的软件并从 NGC 网站复制命令。
然后，通过将命令粘贴到终端中来拉取容器映像。下面的屏幕截图显示了pull命令及其 Docker 操作：

正如您所看到的，Docker 使用基于层的方法。CUDA 容器是在 Ubuntu 的基本层上构建的。此外，Docker images 命令向我们展示了我们机器上本地拉取的容器。

使用以下命令启动拉取的容器：

docker run --rm -it --runtime=nvidia nvcr.io/nvidia/cuda:9.0-devel-ubuntu16.04

GPU 显示在下面的屏幕截图中：

一旦我们运行 Docker，shell 登录会更改，并且我们登录到作为 root 运行的容器中。由于这个原因，我们能够在容器内运行nvidia-smi命令。

我们还可以使用容器通过其附加选项访问主机资源。最常用的选项如下：

- -v：用于挂载卷
-p：用于端口转发
-u：用于用户转发

nvidia-docker的基本用法与普通 Docker 用法类似，只是我们可以使用 GPU。这意味着您还可以获得 Docker 的附加好处。

从 NGC 容器创建和保存新容器

您还可以向现有容器添加层并保存它们以供将来使用。让我们学习如何做到这一点：

创建一个Dockerfile并在基础镜像上创建一些层。例如，我们可以在 NGC PyTorch 容器中更新 APEX（github.com/nvidia/apex）以便我们可以使用其最新版本：

FROM nvcr.io/nvidia/pytorch:19.03-py3
RUN git clone https://github.com/NVIDIA/apex /opt/apex && \
 cd /opt/apex && \
 pip install -v --no-cache-dir --global-option="--cpp_ext" --global-option="--cuda_ext" .

您还可以将所需的 Ubuntu 软件包或 Python 软件包安装代码添加到该文件中。

然后，我们可以使用docker build命令构建一个定制的容器。以下命令显示了 Docker 镜像build命令的基本格式：

docker build -t <image-name>:<tag> .

此命令将找到我们创建的Dockerfile并逐行启动每个命令。Dockerfile的每一行都将创建一个 Docker 层，因此建议编写一个RUN命令来覆盖一个单一目标。

现在，您需要将 Docker 镜像备份到您的私有注册表或创建一个文件。在完成容器后，您可能希望在其他系统中传播或重用该容器。在这种情况下，您可以将 Docker 镜像推送到您的注册表中。例如，如果您在 DockerHub 上有帐户，Docker 提供了一个免费的注册表。您可以使用以下命令将容器推送到注册表中：

docker push <DockerHub-ID>/<image-name>:<tag>

您还可以创建备份文件并将其复制到本地文件系统上。以下命令向您展示了如何使用压缩创建容器备份：

docker save <image-name>:<tag> | gzip > container.tgz

然后，您可以使用以下命令加载该镜像：

gunzip -c container.tgz | docker load

您可以创建一个本地备份镜像而不进行压缩，但通常输出文件太大，无法传送到其他系统。

在本节中，我们已经介绍了 Docker 的一些基本操作。然而，Docker 还提供了其他丰富的功能和好处。尽管 Linux 只能在 Docker 容器中使用 CUDA，但在构建工作环境和帮助您专注于代码开发方面，Docker 会为您节省时间。

将默认运行时设置为 NVIDIA Docker

通过对nvidia-docker配置进行一些修改，我们可以在不通知 GPU 的情况下启动 GPU 容器。因为我们可以将 GPU 运行时选项设置为nvidia-docker，所以我们可以采用 Docker 的运行时设计。为此，您需要将default-runtime": "nvidia",作为选项插入到/etc/docker/daemon.json中。然后，如果没有其他 Docker 配置，可以将daemon.json文件配置如下：

{
    "default-runtime": "nvidia",
    "runtimes": {
        "nvidia": {
            "path": "nvidia-container-runtime",
            "runtimeArgs": []
        }
    }
}

完成此操作后，使用以下命令重新启动系统或重新启动 Docker 守护程序：

sudo systemctl restart docker

现在，我们可以在 Docker 命令中享受 GPU 容器而无需使用 GPU 命令选项。

NVIDIA 开发博客提供了关于nvidia-docker的介绍，可以在devblogs.nvidia.com/gpu-containers-runtime找到。在这里，您不仅将了解其配置，还将了解如何将其与 Docker compose 或 Linux Containers（LXC）集成。它甚至允许 GPU 容器通过其 GPU 设备插件与 Kubernetes 一起工作。

posted @ 2024-05-05 00:03 绝不原创的飞龙阅读(279) 评论(0) 编辑收藏举报

会员力量，点亮园子希望

刷新页面返回顶部

龙哥盟

掠夺·扩张·投机·博弈

CUDA-编程学习手册（全）

CUDA 编程学习手册（全）

前言

这本书适合谁

为了充分利用这本书

下载示例代码文件

下载彩色图像

使用的约定

第一章：CUDA 编程简介

高性能计算的历史

异构计算

编程范式

低延迟与高吞吐量

GPU 的编程方法

技术要求

来自 CUDA 的 Hello World

线程层次结构

GPU 架构

使用 CUDA 进行向量加法

实验 1 - 创建多个块

实验 2 - 创建多个线程

实验 3 - 结合块和线程

为什么要费心处理线程和块？

在多个维度中启动内核

CUDA 中的错误报告

CUDA 中的数据类型支持

总结

第二章：CUDA 内存管理

技术要求

NVIDIA Visual Profiler

全局内存/设备内存

全局内存上的矢量加法

合并与未合并的全局内存访问

内存吞吐量分析

共享内存

共享内存上的矩阵转置

银行冲突及其对共享内存的影响

只读数据/缓存

计算机视觉-使用纹理内存进行图像缩放

GPU 中的寄存器

固定内存

带宽测试-固定与分页

统一内存

了解统一内存页面分配和传输

使用每页 warp 优化统一内存

统一内存的优化使用数据预取

GPU 内存的演变

为什么 GPU 有缓存？

总结

第三章：CUDA 线程编程

技术要求

CUDA 线程、块和 GPU

利用 CUDA 块和 warp

理解 CUDA 占用率

设置 NVCC 报告 GPU 资源使用

Linux 设置

Windows 设置

使用占用率计算器分析最佳占用率

占用率调整 - 限制寄存器使用

从分析器获取实现的占用

理解并行归约

使用全局内存的天真并行归约

使用共享内存减少内核

编写性能测量代码

两种减少-全局和共享内存的性能比较

识别应用程序的性能限制器

找到性能限制因素并进行优化

最小化 CUDA warp 分歧效应

确定分歧作为性能瓶颈

交错寻址

顺序寻址

性能建模和平衡限制器

Roofline 模型

通过网格跨步循环最大化内存带宽

平衡 I/O 吞吐量

warp 级原语编程

使用 warp 原语进行并行归约

灵活处理线程的协作组

使用`grid_group`的用法