Python-GPU-编程实用指南-全-

Python GPU 编程实用指南（全）

原文：zh.annas-archive.org/md5/ef7eb3c148e0cfdfe01c331f2f01557c

译者：飞龙

协议：CC BY-NC-SA 4.0

前言

问候和祝福！本文是关于使用 Python 和 CUDA 进行 GPU 编程的入门指南。GPU可能代表图形编程单元，但我们应该明确，这本书不是关于图形编程——它本质上是通用 GPU 编程的介绍，简称为GPGPU 编程。在过去的十年中，已经清楚地表明 GPU 除了渲染图形之外，也非常适合进行计算，特别是需要大量计算吞吐量的并行计算。为此，NVIDIA 发布了 CUDA 工具包，这使得几乎任何具有一些 C 编程知识的人都可以更轻松地进入 GPGPU 编程的世界。

《使用 Python 和 CUDA 进行 GPU 编程实践》的目标是尽快让您进入 GPGPU 编程的世界。我们努力为每一章设计了有趣和有趣的示例和练习；特别是，我们鼓励您在阅读时在您喜欢的 Python 环境中键入这些示例并运行它们（Spyder、Jupyter 和 PyCharm 都是合适的选择）。这样，您最终将学会所有必要的函数和命令，并获得编写 GPGPU 程序的直觉。

最初，GPGPU 并行编程似乎非常复杂和令人望而生畏，特别是如果您过去只做过 CPU 编程。有很多新概念和约定您必须学习，可能会让您觉得自己从零开始。在这些时候，您必须相信学习这个领域的努力不是徒劳的。通过一点主动性和纪律性，当您阅读完本文时，这个主题将对您来说变得轻而易举。

愉快的编程！

这本书是为谁准备的

这本书特别针对一个人，那就是我自己，2014 年，当时我正在尝试为数学博士研究开发基于 GPU 的模拟。我正在研究多本关于 GPU 编程的书籍和手册，试图对这个领域有一点了解；大多数文本似乎乐意在每一页向读者抛出无尽的硬件原理和术语，而实际的编程则被放在了次要位置。

这本书主要面向那些想要实际进行 GPU 编程，但不想陷入技术细节和硬件原理的人。在本书中，我们将使用适当的 C/C++（CUDA C）来编程 GPU，但我们将通过 PyCUDA 模块将其内联到 Python 代码中。PyCUDA 允许我们只编写我们需要的必要低级 GPU 代码，同时自动处理编译、链接和在 GPU 上启动代码的所有冗余工作。

本书涵盖的内容

第一章《为什么要进行 GPU 编程？》给出了一些我们应该学习这个领域的动机，以及如何应用阿姆达尔定律来估计将串行程序转换为利用 GPU 的潜在性能改进。

第二章《设置 GPU 编程环境》解释了如何在 Windows 和 Linux 下设置适当的 Python 和 C++开发环境以进行 CUDA 编程。

第三章《使用 PyCUDA 入门》展示了我们在使用 Python 编程 GPU 时最基本的技能。我们将特别看到如何使用 PyCUDA 的 gpuarray 类将数据传输到 GPU 和从 GPU 传输数据，以及如何使用 PyCUDA 的 ElementwiseKernel 函数编译简单的 CUDA 核函数。

[第四章]（5a5f4317-50c7-4ce6-9d04-ac3be4c6d28b.xhtml），核心，线程，块和网格，教授了编写有效的 CUDA 核心的基础知识，这些核心是在 GPU 上启动的并行函数。我们将看到如何编写 CUDA 设备函数（由 CUDA 核心直接调用的“串行”函数），并了解 CUDA 的抽象网格/块结构及其在启动核心中的作用。

[第五章]（ea648e20-8c72-44a9-880d-11469d0e291f.xhtml），流，事件，上下文和并发，涵盖了 CUDA 流的概念，这是一种允许我们在 GPU 上同时启动和同步许多内核的功能。我们将看到如何使用 CUDA 事件来计时内核启动，以及如何创建和使用 CUDA 上下文。

[第六章]（6d1c808f-1dc2-4454-b0b8-d0a36bc3c908.xhtml），调试和分析您的 CUDA 代码，填补了我们在纯 CUDA C 编程方面的一些空白，并向我们展示了如何使用 NVIDIA Nsight IDE 进行调试和开发，以及如何使用 NVIDIA 分析工具。

[第七章]（55146879-4b7e-4774-9a8b-cc5c80c04ed8.xhtml），使用 CUDA 库与 Scikit-CUDA，通过 Python Scikit-CUDA 模块简要介绍了一些重要的标准 CUDA 库，包括 cuBLAS，cuFFT 和 cuSOLVER。

[第八章]（d374ea77-f9e5-4d38-861d-5295ef3e3fbf.xhtml），CUDA 设备函数库和 Thrust，向我们展示了如何在我们的代码中使用 cuRAND 和 CUDA Math API 库，以及如何使用 CUDA Thrust C++容器。

[第九章]（3562f1e0-a53d-470f-9b4d-94fa41b1b2fa.xhtml），实现深度神经网络，作为一个巅峰，我们将学习如何从头开始构建整个深度神经网络，应用我们在文本中学到的许多想法。

[第十章]（5383b46f-8dc6-4e17-ab35-7f6bd35f059f.xhtml），使用已编译的 GPU 代码，向我们展示了如何将我们的 Python 代码与预编译的 GPU 代码进行接口，使用 PyCUDA 和 Ctypes。

[第十一章]（e853faad-3ee4-4df7-9cdb-98f74e435527.xhtml），CUDA 性能优化，教授了一些非常低级的性能优化技巧，特别是与 CUDA 相关的技巧，例如 warp shuffling，矢量化内存访问，使用内联 PTX 汇编和原子操作。

[第十二章]（2d464c61-de29-49fa-826a-a7437c368d6a.xhtml），从这里出发，概述了您将拥有的一些教育和职业道路，这些道路将建立在您现在扎实的 GPU 编程基础之上。

为了充分利用本书

这实际上是一个非常技术性的主题。为此，我们将不得不对读者的编程背景做一些假设。为此，我们将假设以下内容：

• 您在 Python 中具有中级的编程经验。

• 您熟悉标准的 Python 科学包，如 NumPy，SciPy 和 Matplotlib。

• 您具有任何基于 C 的编程语言的中级能力（C，C ++，Java，Rust，Go 等）。

• 您了解 C 中的动态内存分配概念（特别是如何使用 C 的malloc和free函数）。

GPU 编程主要适用于非常科学或数学性质的领域，因此许多（如果不是大多数）示例将利用一些数学知识。因此，我们假设读者对大学一年级或二年级的数学有一定了解，包括：

• 三角学（正弦函数：sin，cos，tan…）

• 微积分（积分，导数，梯度）

• 统计学（均匀分布和正态分布）

• 线性代数（向量，矩阵，向量空间，维度）。

如果您还没有学习过这些主题，或者已经有一段时间了，不要担心，因为我们将在学习过程中尝试回顾一些关键的编程和数学概念。

我们在这里将做另一个假设。请记住，本文中我们只会使用 CUDA，这是 NVIDIA 硬件的专有编程语言。因此，在开始之前，我们需要拥有一些特定的硬件。因此，我假设读者可以访问以下内容：

• 64 位 x86 英特尔/AMD PC

• 4 GB 或更多的 RAM

• 入门级 NVIDIA GTX 1050 GPU（Pascal 架构）或更高版本

读者应该知道，大多数旧的 GPU 可能会在本文中的大多数示例中正常工作，但本文中的示例仅在 Windows 10 下的 GTX 1050 和 Linux 下的 GTX 1070 上进行了测试。有关设置和配置的具体说明在第二章中给出，设置您的 GPU 编程环境。

下载示例代码文件

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

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

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

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

3. 单击“代码下载和勘误”。

4. 在搜索框中输入书名，然后按照屏幕上的说明操作。

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

• WinRAR/7-Zip for Windows

• Zipeg/iZip/UnRarX for Mac

• 7-Zip/PeaZip for Linux

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

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

下载彩色图像

我们还提供了一个 PDF 文件，其中包含本书中使用的屏幕截图/图表的彩色图像。您可以在此处下载：www.packtpub.com/sites/default/files/downloads/9781788993913_ColorImages.pdf。

使用的约定

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

CodeInText：表示文本中的代码单词、数据库表名、文件夹名、文件名、文件扩展名、路径名、虚拟 URL、用户输入和 Twitter 句柄。例如：“我们现在可以使用cublasSaxpy函数。”

代码块设置如下：

cublas.cublasDestroy(handle)
print 'cuBLAS returned the correct value: %s' % np.allclose(np.dot(A,x), y_gpu.get())

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

def compute_gflops(precision='S'):

if precision=='S':
    float_type = 'float32'
elif precision=='D':
    float_type = 'float64'
else:
    return -1

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

$ run cublas_gemm_flops.py

粗体：表示新术语、重要单词或屏幕上看到的单词。例如，菜单或对话框中的单词会以这种方式出现在文本中。

警告或重要说明会出现在这样的地方。提示和技巧会出现在这样的地方。

第一章：为什么要进行 GPU 编程？

事实证明，除了能够为视频游戏渲染图形外，图形处理单元（GPU）还为普通消费者提供了一种便捷的方式进行大规模并行 计算——现在普通人可以在当地的电子商店购买一张价值 2000 美元的现代 GPU 卡，将其插入家中的个人电脑，几乎立即就可以用于计算能力，而这种计算能力在 5 年或 10 年前只能在顶级公司和大学的超级计算实验室中获得。近年来，GPU 的开放可访问性在许多方面已经显而易见，这可以通过简要观察新闻来揭示——加密货币挖矿者使用 GPU 生成比特币等数字货币，遗传学家和生物学家使用 GPU 进行 DNA 分析和研究，物理学家和数学家使用 GPU 进行大规模模拟，人工智能研究人员现在可以编程 GPU 来撰写剧本和作曲，而主要的互联网公司，如谷歌和 Facebook，使用带有 GPU 的服务器农场进行大规模机器学习任务……等等。

本书主要旨在让您迅速掌握 GPU 编程，以便您也可以尽快开始使用它们的强大功能，无论您的最终目标是什么。我们旨在涵盖如何编程 GPU 的核心要点，而不是提供 GPU 工作的复杂技术细节和原理图。在本书的末尾，我们将提供更多资源，以便您可以进一步专门化，并应用您对 GPU 的新知识。（有关特定所需的技术知识和硬件的进一步细节，请参阅本节后面的内容。）

在本书中，我们将使用CUDA，这是 NVIDIA 推出的通用 GPU（GPGPU）编程框架，最早发布于 2007 年。虽然 CUDA 专为 NVIDIA GPU 而设计，但它是一个成熟稳定的平台，相对容易使用，提供了一套无与伦比的第一方加速数学和人工智能相关库，并且在安装和集成方面几乎没有麻烦。此外，还有现成的标准化 Python 库，如 PyCUDA 和 Scikit-CUDA，使得渴望成为 GPU 程序员的人更容易接触 GPGPU 编程。出于这些原因，我们选择在本书中使用 CUDA。

CUDA 始终发音为 coo-duh，而不是缩写 C-U-D-A！CUDA 最初代表“计算统一设备架构”，但 Nvidia 已经放弃了这个缩写，现在将 CUDA 作为一个大写的专有名词。

我们现在将开始介绍 GPU 编程的旅程，并概述阿姆达尔定律。阿姆达尔定律是一种简单但有效的方法，用于估计将程序或算法转移到 GPU 上可以获得的潜在速度增益；这将帮助我们确定是否值得重新编写我们的代码以利用 GPU。然后，我们将简要回顾如何使用cProfile模块对我们的 Python 代码进行分析，以帮助我们找到代码中的瓶颈。

本章的学习成果如下：

• 了解阿姆达尔定律

• 在代码的上下文中应用阿姆达尔定律

• 使用cProfile模块对 Python 代码进行基本分析

技术要求

本章建议安装 Anaconda Python 2.7：

www.anaconda.com/download/

本章的代码也可以在 GitHub 上找到：

github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA

有关先决条件的更多信息，请查看本书的前言；有关软件和硬件要求，请查看github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA中的 README 部分。

并行化和阿姆达尔定律

在我们深入了解并发解锁 GPU 的潜力之前，我们首先要意识到它们的计算能力相对于现代英特尔/AMD 中央处理单元（CPU）的优势并不在于它的时钟速度比 CPU 更高，也不在于单个核心的复杂性或特定设计。一个单独的 GPU 核心实际上相当简单，并且与现代单个 CPU 核心相比处于劣势，后者使用了许多花哨的工程技巧，比如分支预测来减少计算的延迟。延迟指的是执行单个计算的开始到结束的持续时间。

GPU 的强大之处在于它的核心比 CPU 多得多，这意味着吞吐量有了巨大的提升。这里的吞吐量指的是可以同时执行的计算数量。让我们使用一个类比来更好地理解这意思。GPU 就像一条非常宽的城市道路，设计成可以同时处理许多行驶缓慢的汽车（高吞吐量，高延迟），而 CPU 就像一条狭窄的高速公路，一次只能容纳几辆车，但可以更快地将每辆车送到目的地（低吞吐量，低延迟）。

我们可以通过查看这些新 GPU 有多少核心来了解吞吐量的增加。举个例子，普通的英特尔或 AMD CPU 只有 2 到 8 个核心，而入门级的消费级 NVIDIA GTX 1050 GPU 有 640 个核心，而新的顶级 NVIDIA RTX 2080 Ti 有 4,352 个核心！我们可以利用这种大规模的吞吐量，只要我们知道如何正确地并行化任何我们希望加速的程序或算法。通过并行化，我们的意思是重写程序或算法，以便我们可以将工作负载并行地在多个处理器上同时运行。让我们从现实生活中想一个类比。

假设你正在建造一座房子，你已经准备好了所有的设计和材料。你雇了一个劳工，你估计需要 100 个小时来建造这座房子。假设这个特定的房子可以以这样的方式建造，即每增加一个劳工，工作就可以完美地分配给他们，也就是说，两个劳工需要 50 个小时，四个劳工需要 25 个小时，十个劳工需要 10 个小时来建造这座房子——建造你的房子所需的时间将是 100 除以你雇佣的劳工数量。这是一个可并行化的任务的例子。

我们注意到，这个任务对于两个劳工来说完成的速度是原来的两倍，对于十个劳工来说完成的速度是原来的十倍（也就是说，并行完成），而不是一个劳工独自建造房子（也就是说，串行完成）——也就是说，如果N是劳工的数量，那么速度将是N倍。在这种情况下，N被称为我们的任务并行化速度的加速比。

在我们开始编写给定算法的并行版本之前，我们经常首先估计一下并行化对我们任务可能带来的潜在 加速。这可以帮助我们确定是否值得花费资源和时间来编写我们程序的并行版本。因为现实生活比我们在这里给出的例子更复杂，很明显我们不可能始终完美地并行化每个程序——大多数情况下，我们的程序只有一部分可以很好地并行化，而其余部分将不得不串行运行。

使用阿姆达尔定律

我们现在将推导阿姆达尔定律，这是一个简单的算术公式，用于估计将一部分串行程序代码并行化到多个处理器上可能带来的潜在速度增益。我们将继续使用我们之前建造房子的类比来做这件事。

上次，我们只考虑了房子的实际物理建造作为整个时间持续时间，但现在，我们还将把设计房子所需的时间考虑在内。假设世界上只有一个人有能力设计你的房子——也就是你——并且你需要 100 小时来设计你的房子的计划。世界上没有其他人能够与你的建筑才华相比，因此这部分任务无法在其他建筑师之间分配，因此无论你拥有什么资源或可以雇佣多少人，设计你的房子都需要 100 小时。因此，如果你只有一名劳工来建造你的房子，建造你的房子所需的整个时间将是 200 小时——你设计它需要 100 小时，一名劳工建造它需要 100 小时。如果我们雇佣两名劳工，这将需要 150 小时——设计房子的时间仍然是 100 小时，而建造将需要 50 小时。很明显，建造房子所需的总时间将是 100 + 100 / N，其中N是我们雇佣的劳工数量。

现在，让我们退一步思考一下，如果我们只雇用一名劳工来建造房子需要多少时间——我们最终使用这个来确定我们雇用额外劳工时的加速度；也就是说，这个过程变得快了多少倍。如果我们只雇用一名劳工，我们会发现设计和建造房子需要相同的时间——100 小时。因此，我们可以说，设计所花费的时间是.5（50%），建造房子所花费的时间也是.5（50%）——当然，这两部分加起来是 1，也就是 100%。当我们增加劳工时，我们想要与这个进行比较——如果我们有两名劳工，建造的时间减半，因此与我们任务的原始串行版本相比，这将花费.5 + .5/2 = .75（75%）的时间，原始任务的.75 x 200 小时是 150 小时，因此我们可以看到这是有效的。此外，我们可以看到，如果我们有N名劳工，我们可以使用公式.5 + .5 / N 来计算我们并行化的建造所需的时间百分比。

现在，让我们确定通过增加额外的劳工我们获得的加速度。如果有两名劳工，建造一座房子只需要 75%的时间，我们可以取.75 的倒数来确定我们并行化的加速度——也就是说，加速度将是 1 / .75，比我们只有一名劳工时快大约 1.33 倍。在这种情况下，我们可以看到，如果有N名劳工，加速度将是 1 / (.5 + .5 / N)。

我们知道，随着我们增加越来越多的劳工，.5 / N 会缩小到接近 0，因此我们可以看到在并行化这个任务时，你可以获得的加速度总是有一个上限，即 1 / (.5 + 0) = 2。我们可以将原始串行时间除以估计的最大加速度，以确定此任务将花费的绝对最短时间——200 / 2 = 100 小时。

我们刚刚应用的用于确定并行编程中加速度的原则被称为阿姆达尔定律。它只需要知道原始串行程序中可并行执行时间的比例，即p，以及我们可用的处理器核心数N。

在这种情况下，不可并行化代码的执行时间比例始终为1-p，因此我们只需要知道p。

我们现在可以使用阿姆达尔定律来计算加速度，如下所示：

总之，Amdahl's Law 是一个简单的公式，允许我们粗略（非常粗略）地估计一个可以至少部分并行化的程序的潜在加速。这可以提供一个大致的想法，即是否值得编写特定串行程序的并行版本，前提是我们知道我们可以并行化代码的比例（p），以及我们可以在其上运行并行化代码的核心数（N）。

Mandelbrot 集

我们现在准备看一个非常标准的并行计算示例，我们将在本文中稍后重新讨论——一个生成Mandelbrot 集图像的算法。让我们首先确切地定义我们的意思。

对于给定的复数c，我们为定义一个递归序列，其中和对于。如果|z[n]|随着n增加到无穷大仍然受到 2 的限制，那么我们将说c是 Mandelbrot 集的成员。

回想一下，我们可以将复数可视化为驻留在二维笛卡尔平面上，其中x轴代表实部分量，y 轴代表虚部分量。因此，我们可以很容易地用一个非常吸引人（并且众所周知）的图形来可视化 Mandelbrot 集。在这里，我们将在复数笛卡尔平面上用浅色表示 Mandelbrot 集的成员，用深色表示非成员，如下所示：

现在，让我们考虑如何在 Python 中生成这个集合。首先，我们必须考虑一些事情——因为显然我们无法检查每一个复数是否在 Mandelbrot 集中，我们必须选择一个特定的范围进行检查；我们必须确定我们将考虑每个范围内的多少点（宽度，高度）；以及我们将检查的|z[n]|的最大值（max_iters）。我们现在可以准备实现一个生成 Mandelbrot 集图形的函数——在这里，我们通过在串行中迭代图中的每一个点来实现这一点。

我们将首先导入 NumPy 库，这是一个我们在本文中将大量使用的数值库。我们在simple_mandelbrot函数中实现这里。我们首先使用 NumPy 的linspace函数生成一个将充当离散复平面的格点（接下来的代码应该相当简单）：

import numpy as np

def simple_mandelbrot(width, height, real_low, real_high, imag_low, imag_high, max_iters):

     real_vals = np.linspace(real_low, real_high, width)
     imag_vals = np.linspace(imag_low, imag_high, height)

     # we will represent members as 1, non-members as 0.

     mandelbrot_graph = np.ones((height,width), dtype=np.float32)

     for x in range(width):

         for y in range(height):

             c = np.complex64( real_vals[x] + imag_vals[y] * 1j  )           
             z = np.complex64(0)

             for i in range(max_iters):

                 z = z**2 + c

                 if(np.abs(z) > 2):
                     mandelbrot_graph[y,x] = 0
                     break

     return mandelbrot_graph

现在，我们想要添加一些代码来将 Mandelbrot 集的图像转储到 PNG 格式文件中，所以让我们在开头添加适当的标头：

from time import time
import matplotlib
# the following will prevent the figure from popping up
matplotlib.use('Agg')
from matplotlib import pyplot as plt

现在，让我们添加一些代码来生成 Mandelbrot 集并将其转储到文件中，并使用时间函数来计算这两个操作的时间：

if __name__ == '__main__':

     t1 = time()
     mandel = simple_mandelbrot(512,512,-2,2,-2,2,256, 2)
     t2 = time()
     mandel_time = t2 - t1

     t1 = time()
     fig = plt.figure(1)
     plt.imshow(mandel, extent=(-2, 2, -2, 2))
     plt.savefig('mandelbrot.png', dpi=fig.dpi)
     t2 = time()

     dump_time = t2 - t1

     print 'It took {} seconds to calculate the Mandelbrot graph.'.format(mandel_time)
     print 'It took {} seconds to dump the image.'.format(dump_time)

现在让我们运行这个程序（这也可以在 GitHub 存储库的文件夹1中的mandelbrot0.py文件中找到）：

生成 Mandelbrot 集大约需要 14.62 秒，转储图像大约需要 0.11 秒。正如我们所看到的，我们逐点生成 Mandelbrot 集；不同点的值之间没有相互依赖，因此，这是一个固有的可并行化函数。相比之下，转储图像的代码无法并行化。

现在，让我们从 Amdahl's Law 的角度来分析这个问题。如果我们在这里并行化我们的代码，我们可以得到什么样的加速？总的来说，程序的两部分共计大约需要 14.73 秒才能运行；因为我们可以并行化 Mandelbrot 集的生成，我们可以说可并行化代码的执行时间部分是 p = 14.62 / 14.73 = .99。这个程序有 99%的可并行性！

我们可能会得到什么样的加速？嗯，我目前正在使用一台配有 640 个核心的入门级 GTX 1050 GPU 的笔记本电脑；因此，当我们使用这个公式时，我们的N将是 640。我们计算速度提升如下：

这绝对非常好，这表明我们值得努力编程使我们的算法使用 GPU。请记住，阿姆达尔定律只是一个非常粗略的估计！当我们将计算卸载到 GPU 时，将会有其他考虑因素，比如 CPU 发送和接收数据到 GPU 的额外时间；或者卸载到 GPU 的算法只能部分并行化。

对代码进行性能分析

在前面的例子中，我们看到我们可以使用 Python 中的标准time函数来分别计时不同的函数和组件。虽然这种方法对我们的小例子程序效果很好，但对于调用许多不同函数的大型程序来说，这种方法并不总是可行，其中一些函数可能值得我们投入精力并行化，或者甚至在 CPU 上进行优化。我们的目标是找到程序的瓶颈和热点，即使我们在每个函数调用周围使用time，我们可能会错过一些东西，或者可能有一些系统或库调用我们甚至没有考虑到，这些调用可能会拖慢速度。在我们考虑重写代码在 GPU 上运行之前，我们必须始终遵循著名的美国计算机科学家唐纳德·克努斯的智慧话语：过早优化是万恶之源。

我们使用所谓的分析器来找到代码中的热点和瓶颈。分析器将方便地让我们看到程序花费最多时间的地方，并让我们相应地进行优化。

使用 cProfile 模块

我们将主要使用cProfile模块来检查我们的代码。这个模块是一个标准库函数，包含在每个现代 Python 安装中。我们可以从命令行运行分析器，使用-m cProfile，并指定我们要按每个函数花费的累积时间来组织结果，然后用>操作符将输出重定向到文本文件中。

这将在 Linux Bash 或 Windows PowerShell 命令行上都可以使用。

现在让我们试一试：

我们现在可以用我们喜欢的文本编辑器查看文本文件的内容。让我们记住，程序的输出将包含在文件的开头：

现在，由于我们没有删除原始示例中对time的引用，我们在开头的前两行看到了它们的输出。然后我们可以看到在程序中进行的函数调用总数，以及运行它所花费的累积时间。

随后，我们有一个程序中调用的函数列表，按照累积耗时最长的函数到最短的顺序排列；第一行是程序本身，而第二行是我们程序中的simple_mandelbrot函数。（请注意，这里的时间与我们用time命令测量的时间一致）。之后，我们可以看到许多与将 Mandelbrot 图形转储到文件相关的库和系统调用，所有这些调用所花费的时间相对较少。我们使用cProfile的输出来推断给定程序中的瓶颈在哪里。

总结

使用 GPU 而不是 CPU 的主要优势是其增加的吞吐量，这意味着我们可以在 GPU 上同时执行更多并行代码，而不是在 CPU 上；GPU 无法使递归算法或非并行化算法变得更快。我们看到一些任务，比如建房子的例子，只能部分并行化——在这个例子中，我们无法加快设计房子的过程（在这种情况下本质上是串行的），但我们可以通过雇佣更多的工人来加快施工的过程（在这种情况下是可并行化的）。

我们使用这个类比来推导阿姆达尔定律，这是一个公式，可以在我们知道可并行代码执行时间百分比以及我们将运行此代码所需的处理器数量时，给出程序潜在加速的粗略估计。然后，我们应用了阿姆达尔定律来分析一个生成曼德勃罗集并将其转储到图像文件的小程序，并且我们确定这将是一个很好的候选者，可以并行化到 GPU 上。最后，我们以简要概述使用cPython模块对代码进行性能分析结束；这使我们能够看到程序中的瓶颈在哪里，而无需显式计时函数调用。

现在我们已经有了一些基本概念，并且有了学习 GPU 编程的动力，我们将在下一章中设置基于 Linux 或 Windows 10 的 GPU 编程环境。然后我们将立即进入 GPU 编程的世界，在接下来的章节中，我们将实际编写一个基于 GPU 的曼德勃罗程序的版本，这是我们在本章中看到的。

问题

1. 本章的曼德勃罗示例中有三个for语句；然而，我们只能并行化前两个。为什么我们不能在所有的for循环上并行化呢？

2. 当我们将阿姆达尔定律应用于将串行 CPU 算法卸载到 GPU 时，有什么是阿姆达尔定律没有考虑到的吗？

3. 假设您独家获得了三个全新的绝密 GPU，它们在所有方面都相同，只是核心数量不同——第一个有 131,072 个核心，第二个有 262,144 个核心，第三个有 524,288 个核心。如果您将曼德勃罗示例并行化并卸载到这些 GPU 上（生成一个 512 x 512 像素的图像），第一个 GPU 和第二个 GPU 之间的计算时间会有区别吗？第二个 GPU 和第三个 GPU 之间呢？

4. 在阿姆达尔定律的背景下，您能想到指定某些算法或代码块为可并行化存在什么问题吗？

5. 为什么我们应该使用性能分析器而不只是使用 Python 的time函数？

第二章：设置 GPU 编程环境

现在我们将看到如何在 Windows 和 Linux 下设置适当的 GPU 编程环境。在这两种情况下，我们都需要采取几个步骤。我们将逐步进行这些步骤，注意 Linux 和 Windows 之间的任何差异。当然，您可以随意跳过或忽略不适用于您选择的操作系统的任何部分或注释。

读者应注意，本章仅涵盖 64 位 Intel/AMD PC 的两个平台——Ubuntu LTS（长期支持）版本和 Windows 10。请注意，任何基于 Ubuntu LTS 的 Linux 操作系统（如 Xubuntu，Kubuntu 或 Linux Mint）也同样适用于通用的 Unity/GNOME-based Ubuntu 版本。

我们建议使用 Python 2.7 而不是 Python 3.x。 Python 2.7 在本文中使用的所有库中都有稳定的支持，并且我们已经在 Windows 和 Linux 平台上使用 Python 2.7 测试了本书中给出的每个示例。 Python 3.x 用户可以使用本书，但应该注意 Python 2.7 和 Python 3.x 之间的区别。本书中的一些示例已经在 Python 3.7 上进行了测试，但需要标准更改，例如在 Python print函数中添加括号。

Packt 作者 Sebastian Raschka 博士在sebastianraschka.com/Articles/2014_python_2_3_key_diff.html提供了 Python 2.7 和 3.x 之间的关键区别列表。

我们特别建议 Windows 和 Linux 用户使用 Anaconda Python 2.7 版本，因为它可以在用户基础上安装，无需sudo或管理员权限，包含本文所需的所有数据科学和可视化模块，并使用快速预优化的 NumPy/SciPy 包，这些包利用了英特尔的数学核心库（MKL）。 （默认的 Linux /usr/bin/python安装对于本文也应该足够，但您可能需要手动安装一些包，如 NumPy 和 Matplotlib。）

Anaconda Python（2.7 和 3.x 版本）可以在www.anaconda.com/download/.下载到所有平台上。

其他受支持平台的用户（例如 macOS，Windows 7/8，Windows Server 2016，Red Hat/Fedora，OpenSUSE 和 CENTOS）应查阅官方的 NVIDIA CUDA 文档（docs.nvidia.com/cuda/）以获取更多详细信息。此外，还有其他硬件选择：对于对嵌入式系统或具有树莓派等开发板经验的读者，可能希望从基于 ARM 的 NVIDIA Jetson 开发板开始，而对于对云计算或 Web 编程感兴趣的读者，可能考虑远程使用适当的 Azure 或 AWS 实例。在这些情况下，鼓励读者阅读官方文档以设置其驱动程序，编译器和 CUDA 工具包。本章中的一些步骤可能适用，也可能不适用。

本章的学习目标是：

• 确保我们拥有适当的硬件

• 安装 NVIDIA GPU 驱动程序

• 设置适当的 C/C++编程环境

• 安装 NVIDIA CUDA 工具包

• 为 GPU 编程设置 Python 环境

技术要求

本章建议安装 Anaconda Python 2.7，网址为www.anaconda.com/download/.

本章的代码也可以在 GitHub 上找到，网址为github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA.

有关先决条件的更多信息，请查看本书的前言；有关软件和硬件要求，请查看github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA中的 README 部分。

确保我们拥有正确的硬件

对于本书，我们建议您至少具备以下硬件：

• 64 位英特尔/AMD PC

• 4GB RAM

• NVIDIA GeForce GTX 1050 GPU（或更高）

这种配置将确保您可以轻松学习 GPU 编程，在本书中运行所有示例，并且还可以运行一些其他新的有趣的基于 GPU 的软件，如 Google 的 TensorFlow（一种机器学习框架）或 Vulkan SDK（一种尖端的图形 API）。

请注意，您必须拥有 NVIDIA 品牌的 GPU 才能使用本书！ CUDA Toolkit 专为 NVIDIA 卡而设计，因此无法用于编程 Intel HD 或 Radeon GPU。

正如所述，我们将假设您使用的是 Windows 10 或 Ubuntu LTS（长期支持）版本。

Ubuntu LTS 发布通常具有 14.04、16.04、18.04 等形式的版本号。

Ubuntu LTS，大体上来说，是最主流的 Linux 版本，可以确保与新软件和工具包的最大兼容性。请记住，有许多基于 Ubuntu 的 Linux 变体，如 Linux Mint 或 Xubuntu，这些通常同样有效。（我个人发现 Linux Mint 在配备 GPU 的笔记本电脑上开箱即用效果相当不错。）

我们应该注意，我们假设您至少拥有一款入门级 GTX 1050（Pascal）GPU，或者在任何更新的架构中具有相当的性能。请注意，本书中的许多示例很可能在大多数旧 GPU 上运行，但作者只在 GTX 1050（在 Windows 10 下）和 GTX 1070（在 Linux 下）上进行了测试。虽然这些示例尚未在旧 GPU 上进行测试，但 2014 年的入门级 Maxwell 架构 GPU，如 GTX 750，也应足以满足本文的要求。

如果您使用的是台式 PC，请确保在继续之前已经按照所有包含的说明物理安装了 GPU。

检查您的硬件（Linux）

现在，我们将在 Linux 中进行一些基本检查，以确保我们拥有正确的硬件。首先让我们打开一个终端并切换到 bash 命令行——您可以通过在 Ubuntu 中快速按下组合键Ctrl + Alt + T来快速完成这一步。

现在，通过输入lscpu并按Enter来检查我们的处理器。会出现大量信息，但只需查看第一行，确保架构确实是 x86_64：

接下来，通过在 bash 提示符下输入free -g并再次按Enter来检查我们的内存容量。这将告诉我们在第一行的第一个条目中我们拥有的总内存量（以 GB 为单位），以及在接下来的行中交换空间中的内存量：

这绝对是足够的内存。

最后，让我们看看我们是否有适当的 GPU。NVIDIA GPU 通过 PCI 总线与我们的 PC 通信，因此我们可以使用lspci命令列出所有 PCI 硬件。通常会列出许多其他硬件，因此让我们使用grep命令仅过滤出 NVIDIA GPU，输入lspci | grep -e "NVIDIA"在 bash 提示符下：

这是一款 GTX 1070，幸运的是它超出了我们至少需要 GTX 1050 的要求。

检查您的硬件（Windows）

首先，我们必须打开 Windows 面板。我们可以通过按下Windows + R，然后在提示符处输入Control Panel来实现这一点，如下面的屏幕截图所示：

Windows 控制面板将弹出。现在点击系统和安全，然后选择以下屏幕上的系统。这将立即告诉我们我们拥有多少 RAM 以及我们是否拥有 64 位处理器：

要检查我们的 GPU，请点击此窗口左上角的设备管理器。然后 Windows 设备管理器将弹出；然后您可以选择显示适配器下拉框来检查您系统上的 GPU：

安装 GPU 驱动程序

如果您已经安装了 GPU 的驱动程序，您可能可以跳过此步骤；此外，一些版本的 CUDA 已经预先打包了最新的驱动程序。通常情况下，CUDA 对您安装的驱动程序非常挑剔，甚至可能无法与 CUDA Toolkit 驱动程序一起工作，因此您可能需要尝试几种不同的驱动程序，直到找到一个可用的。

一般来说，Windows 具有更好的 CUDA 驱动程序兼容性和更用户友好的安装比 Linux。Windows 用户可以考虑跳过此步骤，只使用与 CUDA Toolkit 捆绑的驱动程序，我们稍后将在本章中安装。然而，我们强烈建议 Linux 用户（特别是 Linux 笔记本用户）在继续之前，密切遵循本节中的所有步骤。

安装 GPU 驱动程序（Linux）

在 Ubuntu 中，NVIDIA GPU 的默认驱动程序是一个名为 Nouveau 的开源驱动程序；不幸的是，这在 CUDA 中根本不起作用，因此我们必须安装专有驱动程序。我们必须将特殊的graphics-drivers存储库添加到我们的软件包管理器中，以便能够将专有 NVIDIA 驱动程序下载到我们的 Ubuntu 系统中。我们通过在 bash 提示符中输入以下行来添加存储库：

sudo add-apt-repository ppa:graphics-drivers/ppa

由于这是一个sudo超级用户命令，您将需要输入您的密码。我们现在通过输入以下行来将系统与新的存储库同步：

sudo apt-get update

我们现在应该准备安装我们的驱动程序。从 Ubuntu 桌面，按下Windows + R，然后输入software and drivers：

软件和驱动程序设置菜单应该出现。从这里，点击标记为附加驱动程序的选项卡。您应该看到一系列可用的稳定专有驱动程序供您的 GPU 选择；选择您看到的最新的一个（在我的情况下，它是nvidia-driver-396，如下所示）：

选择最新的驱动程序后，点击应用更改。您将再次被提示输入您的sudo密码，然后驱动程序将安装；进度条应该出现。请注意，这个过程可能需要很长时间，而且可能会出现您的计算机“挂起”的情况；这个过程可能需要超过一个小时，所以请耐心等待。

最后，当过程完成时，重启您的计算机，返回到 Ubuntu 桌面。现在输入Windows + A，然后输入nvidia-settings（或者，从 bash 提示符中运行此程序）。NVIDIA X Server 设置管理器应该出现，并指示您正在使用适当的驱动程序版本：

安装 GPU 驱动程序（Windows）

重申一下-通常建议读者最初跳过此步骤，然后安装包含在 CUDA Toolkit 中的驱动程序。

Windows 的最新驱动程序可以直接从 NVIDIA 的www.nvidia.com/Download/下载。只需从下拉菜单中选择适用于您 GPU 的适当的 Windows 10 驱动程序，这些是可执行（.exe）文件。只需通过双击文件管理器中的文件来安装驱动程序。

建立 C++编程环境

现在我们已经安装了驱动程序，我们必须设置我们的 C/C++编程环境；Python 和 CUDA 都对它们可能集成的编译器和 IDE 有特殊要求，所以您可能需要小心。对于 Ubuntu Linux 用户，标准存储库的编译器和 IDE 通常可以完美地与 CUDA 工具包集成，而 Windows 用户可能需要更加小心。

设置 GCC、Eclipse IDE 和图形依赖项（Linux）

从 Ubuntu 桌面打开终端（Ctrl + Alt + T）。我们首先更新apt存储库如下：

sudo apt-get update

现在我们可以用一行额外的命令安装我们需要的 CUDA 一切：

sudo apt-get install build-essential binutils gdb eclipse-cdt

在这里，build-essential是带有gcc和g++编译器以及其他实用程序（如 make）的软件包；binutils有一些通用的实用程序，如 LD 链接器；gdb是调试器；Eclipse 是我们将要使用的 IDE。

让我们还安装一些额外的依赖项，这将允许我们使用以下命令运行 CUDA 工具包中包含的一些图形（OpenGL）演示：

sudo apt-get install freeglut3 freeglut3-dev libxi-dev libxmu-dev

现在您应该可以安装 CUDA 工具包了。

在 Windows 上设置 Visual Studio

在撰写本文时，只有一个版本的 Visual Studio 似乎完美地与 Python 和最新的 CUDA 工具包集成在一起——Visual Studio 2015；也就是说，Visual Studio 版本 14.0。

虽然可能可以在较新版本的 Visual Studio（例如 2017）下进行子安装，但我们建议读者直接在系统上安装带有 C/C++支持的 Visual Studio 2015。

Visual Studio Community 2015，这个软件的免费版本，可以在visualstudio.microsoft.com/vs/older-downloads/下载。

在这里，我们将进行最小化安装，只安装 CUDA 所需的组件。我们运行安装软件，并选择自定义安装：

点击下一步，然后点击编程语言的下拉框，然后选择 Visual C++（如果您需要其他包或编程语言，可以随意选择，但是对于 GPU 编程，我们只需要 Visual C++）：

这个安装过程可能需要一些时间。完成后，我们将准备安装 CUDA 工具包。

安装 CUDA 工具包

最后，我们开始接近我们的目标！现在我们通过访问developer.nvidia.com/cuda-downloads来下载我们的 CUDA 工具包。选择适当的操作系统，您将看到几个选项。对于 Windows 和 Linux，都有网络和本地安装选项。我倾向于在 Windows 和 Linux 下都使用本地安装选项，因为我更喜欢一次性下载整个软件包；如果有任何网络问题，那么您可以确保在安装 CUDA 工具包时不会发生问题。

安装 CUDA 工具包（Linux）

对于 Linux 用户，您将看到使用.deb包和.run文件的选择；对于大多数用户，我建议使用.deb文件，因为这将自动安装 CUDA 需要的任何缺少的软件包。.run文件安装在系统的高级软件包工具（APT）系统之外，它只是将适当的文件复制到系统的/usr二进制和库目录。如果您不想干扰系统的 APT 系统或存储库，并且对 Linux 有很好的理解，那么.run文件可能更合适。无论哪种情况，请仔细遵循网站上关于安装软件包的说明，这些说明可能会因版本而略有不同。

安装包完成后，您可能需要配置您的PATH和LD_SYSTEM_CONFIG环境变量，以便您的系统可以找到 CUDA 所需的适当的二进制可执行文件和库文件。我建议您通过将以下行附加到您用户目录中的.bashrc文件的末尾来完成这个步骤。使用您喜欢的文本编辑器，如gedit、nano、emacs或vim打开~/.bashrc文件，然后在文件的最底部添加以下行：

export PATH="/usr/local/cuda/bin:${PATH}
export LD_LIBRARY_PATH="/usr/local/cuda/lib64:${LD_LIBRARY_PATH}"

保存文件然后退出终端。您现在可以通过打开一个新的终端并输入nvcc --version然后按Enter来确保您已正确安装了工具包，这将给您工具包编译器的版本信息。（nvcc是命令行 CUDA C 编译器，类似于gcc编译器。）

安装 CUDA Toolkit（Windows）

对于 Windows 用户，您可以通过双击.exe文件并按照屏幕上的提示来安装包。

安装完成后，重置您的系统。我们现在将通过检查nvcc编译器来确保 CUDA 已正确安装。在开始菜单下，点击Visual Studio 2015文件夹，然后点击 VS2015 x64 Native Tools Command Prompt。一个终端窗口将弹出；现在输入nvcc --version并按Enter，这应该会给您 NVIDIA 编译器的版本信息。

为 GPU 编程设置我们的 Python 环境

使用我们的编译器、集成开发环境和 CUDA 工具包正确安装在我们的系统上，我们现在可以为 GPU 编程设置一个合适的 Python 环境。这里有很多选择，但我们明确建议您使用 Anaconda Python Distribution。Anaconda Python 是一个独立且用户友好的分发版，可以直接安装在您的用户目录中，而且不需要任何管理员或sudo级别的系统访问权限来安装、使用或更新。

请记住，Anaconda Python 有两种版本——Python 2.7 和 Python 3。由于 Python 3 目前对我们将要使用的一些库的支持不是很好，我们将在本书中使用 Python 2.7，这仍然是广泛使用的。

您可以通过访问www.anaconda.com/download来安装 Anaconda Python，选择您的操作系统，然后选择下载分发版的 Python 2.7 版本。按照 Anaconda 网站上给出的说明安装分发版，这相对比较简单。现在我们可以为 GPU 编程设置本地 Python 安装。

我们现在将设置本书中可能是最重要的 Python 包：Andreas Kloeckner 的 PyCUDA 包。

安装 PyCUDA（Linux）

在 Linux 中打开一个命令行。通过在 bash 提示符下输入which python并按Enter来确保您的PATH变量正确设置为使用本地 Anaconda 安装的 Python（而不是系统范围的安装）（Anaconda 应该在安装过程中自动配置您的.bashrc）；这应该告诉您 Python 二进制文件在您的本地~/anaconda2/bin目录中，而不是在/usr/bin目录中。如果不是这种情况，请打开一个文本编辑器，并在您的~/.bashrc文件的末尾放置以下行export PATH="/home/${USER}/anaconda2/bin:${PATH}"，保存后，打开一个新的终端，然后再次检查。

有几种安装 PyCUDA 的选项。最简单的选项是从 PyPI 存储库安装最新稳定版本，方法是输入pip install pycuda。您还可以按照 PyCUDA 官方网站上的说明安装最新版本的 PyCUDA，网址为mathema.tician.de/software/pycuda/。请注意，如果您希望从不同的来源重新安装 PyCUDA，请确保首先使用pip uninstall pycuda卸载它。

创建一个环境启动脚本（Windows）

Windows 用户需要特别注意，他们的 Visual Studio 和 Anaconda Python 环境变量是否设置正确，以便使用 PyCUDA；否则，Python 将无法找到 NVIDIA 的nvcc CUDA 编译器或 Microsoft 的cl.exe C++编译器。幸运的是，包含了批处理脚本，可以自动为我们设置这些环境，但我们必须小心，每次想要进行 GPU 编程时都要执行这些脚本。

因此，我们将创建一个批处理脚本，通过连续调用这两个脚本来启动适当的 IDE 或命令行环境。 （此脚本也可在github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA/blob/master/2/launch-python-cuda-environment.bat上找到。）

请务必首先打开 Windows 记事本，并跟随操作：

首先找到您的 Visual Studio 的vcvars.bat文件的位置；对于 Visual Studio 2015，它位于C:\Program Files (x86)\Microsoft Visual Studio 14.0\VC\vcvarsall.bat。

在文本编辑器中输入以下行，然后按Enter：

call "C:\Program Files (x86)\Microsoft Visual Studio 14.0\VC\vcvarsall.bat" amd64

现在，我们需要调用 Anaconda 的activate.bat脚本来设置 Anaconda Python 环境变量；标准路径是Anaconda2\Scripts\activate.bat。我们还必须指示此脚本的参数是 Anaconda 库的位置。在我的情况下，我的启动脚本中的第二行将是call "C:\Users\%username%\Anaconda2\Scripts\activate.bat" C:\Users\%username%\Anaconda2。

最后，我们的批处理脚本的最后一行将启动您喜欢的任何环境——IDE 或命令行提示符，它将继承前两个脚本设置的所有必要环境和系统变量。如果您喜欢旧的标准 DOS 风格命令提示符，这行应该只是cmd。如果您喜欢从 PowerShell 工作，请将其更改为powershell。在某些情况下，特别是用于访问命令行pip和conda来更新 Python 库时，需要使用命令行。

最后，将此文件保存到桌面，并命名为launch-python-cuda-environment.bat。现在，您可以通过双击此文件来启动我们的 Python GPU 编程环境。

（请记住，如果您希望使用 Jupyter Notebook 或 Spyder Python IDE，您可以简单地通过jupyter-notebook或spyder从命令行启动它们，或者您可以制作一个批处理脚本，只需用适当的 IDE 启动命令替换cmd。）

安装 PyCUDA（Windows）

由于大多数 Python 库主要是由 Linux 用户编写和为 Linux 用户编写的，建议您从 Christoph Gohlke 的网站上安装预构建的 PyCUDA wheel 二进制文件，网址为：www.lfd.uci.edu/~gohlke/pythonlibs/#pycuda。下载一个文件，文件名为pycuda-2017.1.1+cuda(VERSION)-cp27-cp27m-win_amd64.whl，其中版本是您的 CUDA 版本号。现在，您可以通过在命令行中输入以下内容并用您的 PyCUDA wheel 的完整路径和文件名替换pycuda.whl来安装 PyCUDA：

pip install pycuda.whl

（或者，您可以尝试使用pip install pycuda从 PyPI 存储库安装 PyCUDA，或者按照 PyCUDA 网站上的说明操作。）

测试 PyCUDA

最后，我们到了可以看到我们的 GPU 编程环境是否真正起作用的时候。我们将运行下一章的一个小程序，该程序将查询我们的 GPU 并提供有关型号号码、内存、核心数量、架构等相关信息。从存储库中的目录3中获取 Python 文件（deviceQuery.py），也可以在github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA/blob/master/3/deviceQuery.py上找到。

如果您使用的是 Windows，请确保通过在桌面上启动我们在上一节中创建的.bat文件来启动 GPU 编程环境。否则，如果您使用的是 Linux，请打开一个 bash 终端。现在输入以下命令并按Enter键——python deviceQuery.py。

这将输出许多行数据，但前几行应该表明 PyCUDA 已经检测到您的 GPU，并且您应该在下一行看到型号号码：

恭喜，您现在已经准备好进入 GPU 编程的世界了！

总结

为 GPU 编程设置 Python 环境可能是一个非常微妙的过程。本文建议 Windows 和 Linux 用户都使用 Anaconda Python 2.7 发行版。首先，我们应该确保我们有正确的硬件进行 GPU 编程；一般来说，64 位 Windows 或 Linux PC，带有 4GB RAM 和 2016 年或之后的任何入门级 NVIDIA GPU 将足够满足我们的需求。Windows 用户应该注意使用一个既适用于 CUDA 工具包又适用于 Anaconda 的 Visual Studio 版本（如 VS 2015），而 Linux 用户在安装 GPU 驱动程序时应特别小心，并在其.bashrc文件中设置适当的环境变量。此外，Windows 用户应该创建一个适当的启动脚本，用于设置 GPU 编程环境，并应该使用预编译的 PyCUDA 库安装文件。

现在，我们的编程环境已经设置好了，接下来的一章我们将学习 GPU 编程的基础知识。我们将看到如何将数据写入 GPU 的内存，以及如何在 CUDA C 中编写一些非常简单的逐元素GPU 函数。（如果你看过经典的 1980 年代电影《功夫小子》，那么你可能会把下一章看作是 GPU 编程的“上蜡，下蜡”。）

问题

1. 我们可以在主处理器内置的英特尔 HD GPU 上运行 CUDA 吗？离散的 AMD Radeon GPU 呢？

2. 这本书的示例是使用 Python 2.7 还是 Python 3.7？

3. 在 Windows 中，我们使用什么程序来查看我们安装了什么 GPU 硬件？

4. 在 Linux 中，我们使用什么命令行程序来查看我们安装了什么 GPU 硬件？

5. 在 Linux 中，我们使用什么命令来确定系统有多少内存？

6. 如果我们不想改变我们的 Linux 系统的 APT 存储库，我们应该使用run还是deb安装程序来安装 CUDA？

第三章：开始使用 PyCUDA

在上一章中，我们设置了编程环境。现在，有了我们的驱动程序和编译器牢固地安装好，我们将开始实际的 GPU 编程！我们将首先学习如何使用 PyCUDA 进行一些基本和基础的操作。我们将首先看看如何查询我们的 GPU - 也就是说，我们将首先编写一个小的 Python 程序，告诉我们 GPU 的特性，如核心数量、架构和内存。然后，我们将花一些时间熟悉如何在 Python 和 GPU 之间传输内存，使用 PyCUDA 的gpuarray类以及如何使用这个类进行基本计算。本章的其余部分将花在展示如何编写一些基本函数（我们将称之为CUDA 内核），我们可以直接启动到 GPU 上。

本章的学习成果如下：

• 使用 PyCUDA 确定 GPU 特性，如内存容量或核心数量

• 理解主机（CPU）和设备（GPU）内存之间的区别，以及如何使用 PyCUDA 的gpuarray类在主机和设备之间传输数据

• 如何使用只有gpuarray对象进行基本计算

• 如何使用 PyCUDA 的ElementwiseKernel函数在 GPU 上执行基本的逐元素操作

• 理解函数式编程概念的 reduce/scan 操作，以及如何制作基本的缩减或扫描 CUDA 内核

技术要求

本章需要一台安装了现代 NVIDIA GPU（2016 年以后）的 Linux 或 Windows 10 PC，并安装了所有必要的 GPU 驱动程序和 CUDA Toolkit（9.0 以后）。还需要一个合适的 Python 2.7 安装（如 Anaconda Python 2.7），并安装了 PyCUDA 模块。

本章的代码也可以在 GitHub 上找到：github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA。

有关先决条件的更多信息，请查看本书的前言；有关软件和硬件要求，请查看github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA中的README部分。

查询您的 GPU

在我们开始编程 GPU 之前，我们应该真正了解一些关于其技术能力和限制的知识。我们可以通过进行所谓的GPU 查询来确定这一点。GPU 查询是一个非常基本的操作，它将告诉我们 GPU 的具体技术细节，如可用的 GPU 内存和核心数量。NVIDIA 在samples目录中（适用于 Windows 和 Linux）包含了一个纯 CUDA-C 编写的命令行示例deviceQuery，我们可以运行它来执行此操作。让我们看一下作者的 Windows 10 笔记本电脑（Microsoft Surface Book 2，配备了 GTX 1050 GPU）上产生的输出：

让我们来看看这里显示的所有技术信息的一些基本要点。首先，我们看到只安装了一个 GPU，设备 0 - 可能主机计算机安装了多个 GPU 并使用它们，因此 CUDA 将为每个GPU 设备指定一个独立的编号。有些情况下，我们可能需要明确指定设备编号，所以了解这一点总是很好的。我们还可以看到我们拥有的具体设备类型（这里是 GTX 1050），以及我们正在使用的 CUDA 版本。现在我们还要注意两件事：核心的总数（这里是 640），以及设备上的全局内存总量（在本例中为 2,048 兆字节，即 2 千兆字节）。

虽然您可以从deviceQuery中看到许多其他技术细节，但核心数量和内存量通常是您第一次在新 GPU 上运行时应该关注的前两件事，因为它们可以让您最直接地了解新设备的容量。

使用 PyCUDA 查询您的 GPU

现在，最后，我们将通过用 Python 编写我们自己的版本的deviceQuery来开始我们的 GPU 编程之旅。在这里，我们主要关注设备上可用内存的数量，计算能力，多处理器的数量和 CUDA 核心的总数。

我们将从以下方式初始化 CUDA：

import pycuda.driver as drv
drv.init()

请注意，我们将始终需要使用pycuda.driver.init()或通过导入 PyCUDA 的autoinit子模块import pycuda.autoinit来初始化 PyCUDA！

我们现在可以立即检查我们的主机计算机上有多少个 GPU 设备：

print 'Detected {} CUDA Capable device(s)'.format(drv.Device.count())

让我们在 IPython 中输入这个并看看会发生什么：

太好了！到目前为止，我已经验证了我的笔记本确实有一个 GPU。现在，让我们通过添加几行代码来迭代可以通过pycuda.driver.Device（按编号索引）单独访问的每个设备，以提取有关此 GPU（以及系统上的任何其他 GPU）的更多有趣信息。设备的名称（例如，GeForce GTX 1050）由name函数给出。然后我们使用compute_capability函数获取设备的计算能力和total_memory函数获取设备的总内存量。

计算能力可以被视为每个 NVIDIA GPU 架构的版本号；这将为我们提供一些关于设备的重要信息，否则我们无法查询，我们将在一分钟内看到。

我们将这样写：

for i in range(drv.Device.count()):

     gpu_device = drv.Device(i)
     print 'Device {}: {}'.format( i, gpu_device.name() )
     compute_capability = float( '%d.%d' % gpu_device.compute_capability() )
     print '\t Compute Capability: {}'.format(compute_capability)
     print '\t Total Memory: {} megabytes'.format(gpu_device.total_memory()//(1024**2))

现在，我们准备查看 PyCUDA 以 Python 字典类型形式提供给我们的 GPU 的一些剩余属性。我们将使用以下行将其转换为由字符串索引属性的字典：

    device_attributes_tuples = gpu_device.get_attributes().iteritems()
     device_attributes = {}

     for k, v in device_attributes_tuples:
         device_attributes[str(k)] = v

现在，我们可以使用以下内容确定设备上的多处理器数量：

    num_mp = device_attributes['MULTIPROCESSOR_COUNT']

GPU 将其各个核心划分为称为流处理器（SMs）的较大单元； GPU 设备将具有多个 SM，每个 SM 将根据设备的计算能力具有特定数量的 CUDA 核心。要明确：每个多处理器的核心数并不是由 GPU 直接指示的-这是由计算能力隐含给我们的。我们将不得不查阅 NVIDIA 的一些技术文件以确定每个多处理器的核心数（参见docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capabilities），然后创建一个查找表来给出每个多处理器的核心数。我们使用compute_capability变量来查找核心数：

    cuda_cores_per_mp = { 5.0 : 128, 5.1 : 128, 5.2 : 128, 6.0 : 64, 6.1 : 128, 6.2 : 128}[compute_capability]

现在我们可以通过将这两个数字相乘来最终确定设备上的总核心数：

    print '\t ({}) Multiprocessors, ({}) CUDA Cores / Multiprocessor: {} CUDA Cores'.format(num_mp, cuda_cores_per_mp, num_mp*cuda_cores_per_mp)

现在，我们可以通过迭代字典中剩余的键并打印相应的值来完成我们的程序：

    device_attributes.pop('MULTIPROCESSOR_COUNT')

     for k in device_attributes.keys():
         print '\t {}: {}'.format(k, device_attributes[k])

所以，现在我们终于完成了本文的第一个真正的 GPU 程序！（也可在github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA/blob/master/3/deviceQuery.py找到）。现在，我们可以按如下方式运行它：

现在，我们可以有点自豪，因为我们确实可以编写一个程序来查询我们的 GPU！现在，让我们真正开始学习使用我们的 GPU，而不仅仅是观察它。

使用 PyCUDA 的 gpuarray 类

就像 NumPy 的 array 类是 NumPy 环境中数值编程的基石一样，PyCUDA 的 gpuarray 类在 Python 中的 GPU 编程中扮演着类似的重要角色。它具有你从 NumPy 中熟悉和喜爱的所有功能——多维向量/矩阵/张量形状结构、数组切片、数组展开，以及用于逐点计算的重载运算符（例如 +、-、*、/ 和 **）。

gpuarray 对于任何新手 GPU 程序员来说都是一个不可或缺的工具。在我们继续之前，我们将花费这一部分时间来了解这种特定的数据结构，并对其有一个深入的理解。

使用 gpuarray 将数据传输到 GPU 和从 GPU 中传输数据

正如我们在 Python 中编写的先前的 deviceQuery 程序所示，GPU 有自己的内存，与主机计算机的内存分开，这被称为设备内存。（有时这更具体地被称为全局设备内存，以区分它与 GPU 上的其他缓存内存、共享内存和寄存器内存。）在大多数情况下，我们将 GPU 上的（全局）设备内存视为我们在 C 中动态分配的堆内存（使用 malloc 和 free 函数）或 C++（使用 new 和 delete 运算符）；在 CUDA C 中，这进一步复杂化，需要在 CPU 和 GPU 空间之间来回传输数据（使用诸如 cudaMemcpyHostToDevice 和 cudaMemcpyDeviceToHost 的命令），同时跟踪 CPU 和 GPU 空间中的多个指针，并执行适当的内存分配（cudaMalloc）和释放（cudaFree）。

幸运的是，PyCUDA 通过 gpuarray 类涵盖了所有的内存分配、释放和数据传输的开销。正如所述，这个类类似于 NumPy 数组，使用矢量/矩阵/张量形状结构信息来处理数据。gpuarray 对象甚至根据其生命周期自动执行清理，因此当我们完成后，我们不必担心释放存储在 gpuarray 对象中的任何 GPU 内存。

那么我们如何使用它将数据从主机传输到 GPU 呢？首先，我们必须将我们的主机数据包含在某种形式的 NumPy 数组中（我们称之为 host_data），然后使用 gpuarray.to_gpu(host_data) 命令将其传输到 GPU 并创建一个新的 GPU 数组。

现在让我们在 GPU 中执行一个简单的计算（在 GPU 上的常数点乘），然后使用 gpuarray.get 函数将 GPU 数据检索到一个新的数组中。让我们加载 IPython 并看看它是如何工作的（请注意，这里我们将使用 import pycuda.autoinit 初始化 PyCUDA）：

需要注意的一点是，当我们设置 NumPy 数组时，我们特别指定了主机上的数组类型为 NumPy float32 类型，并使用 dtype 选项；这与 C/C++ 中的浮点类型直接对应。一般来说，当我们发送数据到 GPU 时，最好使用 NumPy 明确设置数据类型。原因有两个：首先，由于我们使用 GPU 来提高应用程序的性能，我们不希望使用不必要的类型造成不必要的计算时间或内存开销；其次，由于我们很快将在内联 CUDA C 中编写代码的部分，我们必须非常具体地指定类型，否则我们的代码将无法正确工作，要记住 C 是一种静态类型语言。

记得为将要传输到 GPU 的 NumPy 数组明确设置数据类型。这可以在 numpy.array 类的构造函数中使用 dtype 选项来完成。

使用 gpuarray 进行基本的逐点算术运算

在最后一个示例中，我们看到我们可以使用（重载的）Python 乘法运算符（*）来将gpuarray对象中的每个元素乘以一个标量值（这里是 2）；请注意，逐点操作本质上是可并行化的，因此当我们在gpuarray对象上使用此操作时，PyCUDA 能够将每个乘法操作分配到单个线程上，而不是依次串行计算每个乘法（公平地说，一些版本的 NumPy 可以使用现代 x86 芯片中的高级 SSE 指令进行这些计算，因此在某些情况下性能将与 GPU 相当）。明确一点：在 GPU 上执行的这些逐点操作是并行的，因为一个元素的计算不依赖于任何其他元素的计算。

为了了解这些运算符的工作原理，我建议读者加载 IPython 并在 GPU 上创建一些gpuarray对象，然后玩几分钟，看看这些运算符是否与 NumPy 中的数组类似。以下是一些灵感：

现在，我们可以看到gpuarray对象的行为是可预测的，并且与 NumPy 数组的行为一致。（请注意，我们将不得不使用get函数从 GPU 中获取输出！）现在让我们比较一下 CPU 和 GPU 计算时间，看看在何时是否有任何优势进行这些操作。

速度测试

让我们编写一个小程序（time_calc0.py），对 CPU 上的标量乘法和 GPU 上的相同操作进行速度比较测试。然后，我们将使用 NumPy 的allclose函数比较两个输出值。我们将生成一个包含 5000 万个随机 32 位浮点值的数组（这将大约占用 48 兆字节的数据，因此在任何稍微现代的主机和 GPU 设备上都应该完全可行），然后我们将计算在两个设备上将数组乘以 2 所需的时间。最后，我们将比较输出值以确保它们相等。操作如下：

import numpy as np
import pycuda.autoinit
from pycuda import gpuarray
from time import time
host_data = np.float32( np.random.random(50000000) )

t1 = time()
host_data_2x =  host_data * np.float32(2)
t2 = time()

print 'total time to compute on CPU: %f' % (t2 - t1)
device_data = gpuarray.to_gpu(host_data)

t1 = time()
device_data_2x =  device_data * np.float32( 2 )
t2 = time()

from_device = device_data_2x.get()
print 'total time to compute on GPU: %f' % (t2 - t1)

print 'Is the host computation the same as the GPU computation? : {}'.format(np.allclose(from_device, host_data_2x) )

（您可以在之前提供给您的存储库中找到time_calc0.py文件。）

现在，让我们加载 IPython 并运行几次，以了解这些的一般速度，并查看是否有任何变化。（这里，这是在 2017 年的微软 Surface Book 2 上运行的，配备了 Kaby Lake i7 处理器和 GTX 1050 GPU。）：

我们首先注意到，每次计算的 CPU 计算时间大致相同（大约 0.08 秒）。然而，我们注意到，第一次运行时，GPU 计算时间比 CPU 计算时间慢得多（1.09 秒），并且在随后的运行中变得快得多，在每次后续运行中保持大致恒定（在 7 或 9 毫秒的范围内）。如果您退出 IPython，然后再次运行程序，将发生相同的情况。这种现象的原因是什么？好吧，让我们使用 IPython 内置的prun分析器进行一些调查工作。（这类似于第一章中介绍的cProfiler模块，为什么要进行 GPU 编程？）

首先，让我们将我们的程序作为文本加载到 IPython 中，然后通过 Python 的exec命令运行我们的分析器：

with open('time_calc0.py','r') as f:
     time_calc_code = f.read()

现在，我们在 IPython 控制台中键入%prun -s cumulative exec(time_calc_code)（带有前导%）并查看哪些操作花费了最多的时间：

现在，有一些可疑的对 Python 模块文件compiler.py的调用；这些调用总共大约需要一秒钟，比在这里进行 GPU 计算所需的时间略少。现在让我们再次运行一下，看看是否有任何差异：

请注意，这一次没有调用compiler.py。为什么呢？由于 PyCUDA 库的性质，GPU 代码通常在给定的 Python 会话中首次运行时使用 NVIDIA 的nvcc编译器进行编译和链接；然后它被缓存，如果再次调用代码，则不必重新编译。这甚至可能包括简单的操作，比如标量乘法！（我们最终会看到，通过使用第十章中的预编译代码或使用 NVIDIA 自己的线性代数库与 Scikit-CUDA 模块一起使用 CUDA 库，可以改善这一点，我们将在第七章中看到）。

在 PyCUDA 中，GPU 代码通常在运行时使用 NVIDIA 的nvcc编译器进行编译，然后从 PyCUDA 中调用。这可能会导致意外的减速，通常是在给定的 Python 会话中首次运行程序或 GPU 操作时。

使用 PyCUDA 的 ElementWiseKernel 执行逐点计算

现在，让我们看看如何使用 PyCUDA 的ElementWiseKernel函数直接在 GPU 上编写我们自己的逐点（或等效地，逐元素）操作。这就是我们之前对 C/C++编程的了解将变得有用的地方——我们将不得不在 CUDA C 中编写一点内联代码，这些代码是由 NVIDIA 的nvcc编译器在外部编译的，然后通过 PyCUDA 在运行时由我们的代码启动。

在本文中，我们经常使用术语kernel；通过kernel，我们总是指的是由 CUDA 直接启动到 GPU 上的函数。我们将使用 PyCUDA 的几个函数来生成不同类型的 kernel 的模板和设计模式，以便更轻松地过渡到 GPU 编程。

让我们直接开始；我们将从头开始重写代码，使用 CUDA-C 将gpuarray对象的每个元素乘以 2；我们将使用 PyCUDA 的ElementwiseKernel函数来生成我们的代码。您应该尝试直接在 IPython 控制台中输入以下代码。（不那么冒险的人可以从本文的 Git 存储库中下载，文件名为simple_element_kernel_example0.py）：

import numpy as np
import pycuda.autoinit
from pycuda import gpuarray
from time import time
from pycuda.elementwise import ElementwiseKernel
host_data = np.float32( np.random.random(50000000) )
gpu_2x_ker = ElementwiseKernel(
"float *in, float *out",
"out[i] = 2*in[i];",
"gpu_2x_ker")

让我们看看这是如何设置的；当然，这是几行内联 C。我们首先在第一行中设置输入和输出变量（"float *in, float *out"），这通常是指指向 GPU 上已分配内存的 C 指针的形式。在第二行中，我们使用"out[i] = 2*in[i];"定义了我们的逐元素操作，它将把in中的每个点乘以 2，并将其放在out的相应索引中。

请注意，PyCUDA 会自动为我们设置整数索引i。当我们使用i作为我们的索引时，ElementwiseKernel将自动在 GPU 的许多核心中并行化我们的计算。最后，我们给我们的代码片段起了一个内部 CUDA C kernel 的名称（"gpu_2x_ker"）。由于这是指 CUDA C 的命名空间而不是 Python 的，因此将其与 Python 中的名称相同是可以的（也很方便）。

现在，让我们进行速度比较：

def speedcomparison():
    t1 = time()
    host_data_2x =  host_data * np.float32(2)
    t2 = time()
    print 'total time to compute on CPU: %f' % (t2 - t1)
    device_data = gpuarray.to_gpu(host_data)
    # allocate memory for output
    device_data_2x = gpuarray.empty_like(device_data)
    t1 = time()
    gpu_2x_ker(device_data, device_data_2x)
    t2 = time()
    from_device = device_data_2x.get()
    print 'total time to compute on GPU: %f' % (t2 - t1)
    print 'Is the host computation the same as the GPU computation? : {}'.format(np.allclose(from_device, host_data_2x) )

if __name__ == '__main__':
    speedcomparison()

现在，让我们运行这个程序：

哇！看起来不太好。让我们从 IPython 中运行speedcomparison()函数几次：

正如我们所看到的，第一次使用给定的 GPU 函数后，速度显著增加。与前面的例子一样，这是因为 PyCUDA 在首次调用给定的 GPU kernel 函数时使用nvcc编译器编译我们的内联 CUDA C 代码。代码编译后，它将被缓存并在给定的 Python 会话的其余部分中重复使用。

现在，在我们继续之前，让我们再讨论一些重要的事情，这是非常微妙的。我们定义的小内核函数操作 C 浮点指针；这意味着我们将不得不在 GPU 上分配一些空的内存，该内存由out变量指向。再次看一下speedcomparison()函数中的这部分代码：

device_data = gpuarray.to_gpu(host_data)
# allocate memory for output
device_data_2x = gpuarray.empty_like(device_data)

与之前一样，我们通过gpuarray.to_gpu函数将一个 NumPy 数组（host_data）发送到 GPU，该函数会自动将数据分配到 GPU 并从 CPU 空间复制过来。我们将把这个数组插入到我们内核函数的in部分。在下一行，我们使用gpuarray.empty_like函数在 GPU 上分配空的内存。这类似于 C 中的普通malloc，分配一个与device_data大小和数据类型相同的数组，但不复制任何内容。现在我们可以将其用于内核函数的out部分。现在我们来看一下speedcomparison()中的下一行，看看如何将我们的内核函数启动到 GPU 上（忽略我们用于计时的行）：

gpu_2x_ker(device_data, device_data_2x)

再次，我们设置的变量直接对应于我们用ElementwiseKernel定义的第一行（这里是"float *in, float *out"）。

曼德勃罗重新审视

让我们再次从第一章“为什么使用 GPU 编程？”中生成曼德勃罗集的问题。原始代码可以在存储库的1文件夹中找到，文件名为mandelbrot0.py，在继续之前，您应该再次查看一下。我们看到该程序有两个主要组成部分：第一个是生成曼德勃罗集，第二个是将曼德勃罗集转储到 PNG 文件中。在第一章中，我们意识到我们只能并行生成曼德勃罗集，并且考虑到这占程序运行时间的大部分，这将是一个很好的候选算法，可以将其转移到 GPU 上。让我们看看如何做到这一点。（我们将避免重复定义曼德勃罗集，因此如果您需要更深入的复习，请重新阅读第一章“为什么使用 GPU 编程？”中的“曼德勃罗重新审视”部分）

首先，让我们基于原始程序中的simple_mandelbrot创建一个新的 Python 函数。我们将其称为gpu_mandelbrot，这将接受与之前完全相同的输入：

def gpu_mandelbrot(width, height, real_low, real_high, imag_low, imag_high, max_iters, upper_bound):

我们将从这里开始以稍微不同的方式进行。我们将首先构建一个复杂的晶格，其中包含我们将分析的复平面中的每个点。

在这里，我们将使用一些 NumPy 矩阵类型的技巧轻松生成晶格，然后将结果从 NumPy matrix类型转换为二维 NumPy array类型（因为 PyCUDA 只能处理 NumPy array类型，而不能处理matrix类型）。请注意我们非常小心地设置我们的 NumPy 类型：

    real_vals = np.matrix(np.linspace(real_low, real_high, width), dtype=np.complex64)
    imag_vals = np.matrix(np.linspace( imag_high, imag_low, height), dtype=np.complex64) * 1j
    mandelbrot_lattice = np.array(real_vals + imag_vals.transpose(), dtype=np.complex64)  

因此，我们现在有一个表示我们将生成曼德勃罗集的晶格的二维复杂数组；正如我们将看到的，我们可以在 GPU 内非常容易地操作这个数组。现在让我们将我们的晶格传输到 GPU，并分配一个数组来表示我们的曼德勃罗集：

    # copy complex lattice to the GPU
    mandelbrot_lattice_gpu = gpuarray.to_gpu(mandelbrot_lattice)    
    # allocate an empty array on the GPU
    mandelbrot_graph_gpu = gpuarray.empty(shape=mandelbrot_lattice.shape, dtype=np.float32)

重申一下——gpuarray.to_array函数只能操作 NumPy array类型，因此我们在将其发送到 GPU 之前一定要对其进行类型转换。接下来，我们必须使用gpuarray.empty函数在 GPU 上分配一些内存，指定数组的大小/形状和类型。同样，您可以将其视为类似于 C 中的malloc；请记住，由于gpuarray对象析构函数在作用域结束时自动处理内存清理，因此我们不必在以后释放或free这些内存。

当您使用 PyCUDA 函数gpuarray.empty或gpuarray.empty_like在 GPU 上分配内存时，由于gpuarray对象的析构函数管理所有内存清理，因此您不必在以后释放此内存。

现在我们准备启动内核；我们唯一需要做的更改是改变

我们还没有编写生成曼德勃罗集的内核函数，但让我们先写出这个函数的其余部分应该是怎样的：

    mandel_ker( mandelbrot_lattice_gpu, mandelbrot_graph_gpu, np.int32(max_iters), np.float32(upper_bound))

    mandelbrot_graph = mandelbrot_graph_gpu.get()

    return mandelbrot_graph

这就是我们希望我们的新内核行为的方式——第一个输入将是我们生成的复点阵（NumPy complex64类型），第二个将是指向二维浮点数组的指针（NumPy float32类型），它将指示哪些元素是曼德勃罗集的成员，第三个将是一个整数，表示每个点的最大迭代次数，最后一个输入将是用于确定曼德勃罗类成员资格的每个点的上限。请注意，我们在将所有输入传递给 GPU 时非常小心！

下一行将从 GPU 中检索我们生成的曼德勃罗集回到 CPU 空间，并返回结束值。（请注意，gpu_mandelbrot的输入和输出与simple_mandelbrot完全相同）。

现在让我们看看如何正确定义我们的 GPU 内核。首先，让我们在头部添加适当的include语句：

import pycuda.autoinit
from pycuda import gpuarray
from pycuda.elementwise import ElementwiseKernel

我们现在准备编写我们的 GPU 内核！我们将在这里展示它，然后逐行讨论：

mandel_ker = ElementwiseKernel(
"pycuda::complex<float> *lattice, float *mandelbrot_graph, int max_iters, float upper_bound",
"""
mandelbrot_graph[i] = 1;
pycuda::complex<float> c = lattice[i]; 
pycuda::complex<float> z(0,0);
for (int j = 0; j < max_iters; j++)
    {  
     z = z*z + c;     
     if(abs(z) > upper_bound)
         {
          mandelbrot_graph[i] = 0;
          break;
         }
    }         
""",
"mandel_ker")

首先，我们使用传递给ElementwiseKernel的第一个字符串设置我们的输入。我们必须意识到当我们在 CUDA-C 中工作时，特定的 C 数据类型将直接对应于特定的 Python NumPy 数据类型。再次注意，当数组被传递到 CUDA 内核时，它们被 CUDA 视为 C 指针。在这里，CUDA C int类型与 NumPy int32类型完全对应，而 CUDA C float类型对应于 NumPy float32类型。然后使用内部 PyCUDA 类模板进行复杂类型的转换——这里 PyCUDA ::complex<float>对应于 Numpy complex64。

让我们看看第二个字符串的内容，它用三个引号（"""）分隔。这使我们能够在字符串中使用多行；当我们在 Python 中编写更大的内联 CUDA 内核时，我们将使用这个。

虽然我们传入的数组在 Python 中是二维数组，但 CUDA 只会将它们视为一维数组，并由i索引。同样，ElementwiseKernel会自动为我们跨多个核心和线程索引i。我们将输出中的每个点初始化为 1，如mandelbrot_graph[i] = 1;，因为i将在曼德勃罗集的每个元素上进行索引；我们将假设每个点都是成员，除非证明相反。（再次说明，曼德勃罗集是在两个维度上的，实部和虚部，但ElementwiseKernel将自动将所有内容转换为一维集合。当我们再次在 Python 中与数据交互时，曼德勃罗集的二维结构将被保留。）

我们像在 Python 中一样为我们的c值设置适当的点阵点，如pycuda::complex<float> c = lattice[i];，并用pycuda::complex<float> z(0,0);将我们的z值初始化为0（第一个零对应于实部，而第二个对应于虚部）。然后我们使用一个新的迭代器j进行循环，如for(int j = 0; j < max_iters; j++)。（请注意，这个算法不会在j或任何其他索引上并行化——只有i！这个for循环将在j上串行运行——但整个代码片段将在i上并行化。）

然后，我们使用z = z*z + c;设置*z*的新值，按照曼德勃罗算法。如果这个元素的绝对值超过了上限（if(abs(z) > upper_bound)），我们将这个点设置为 0（mandelbrot_graph[i] = 0;），并用break关键字跳出循环。

在传递给ElementwiseKernel的最终字符串中，我们为内核赋予其内部 CUDA C 名称，这里是"mandel_ker"。

我们现在准备启动内核；我们唯一需要做的更改是将主函数中的引用从simple_mandelbrot更改为gpu_mandelbrot，然后我们就可以开始了。让我们从 IPython 中启动：

让我们检查转储的图像，以确保这是正确的：

这肯定是在第一章中生成的相同 Mandelbrot 图像，所以我们已经成功地将其实现到了 GPU 上！现在让我们看看我们得到的速度增加：在第一章中，我们花了 14.61 秒来生成这张图；而在这里，只花了 0.894 秒。请记住，PyCUDA 还必须在运行时编译和链接我们的 CUDA C 代码，并且需要花费时间来进行与 GPU 的内存传输。即使有了所有这些额外的开销，它仍然是一个非常值得的速度增加！（您可以在 Git 存储库中找到我们的 GPU Mandelbrot 的代码，文件名为gpu_mandelbrot0.py。）

对函数式编程的简要探讨

在我们继续之前，让我们简要回顾一下 Python 中用于函数式编程的两个函数——map和reduce。它们都被认为是函数式，因为它们都对函数进行操作。我们发现这些有趣，因为它们都对应于编程中的常见设计模式，所以我们可以替换输入中的不同函数，以获得多种不同（和有用的）操作。

让我们首先回顾 Python 中的lambda关键字。这允许我们定义一个匿名函数——在大多数情况下，这些可以被视为一次性函数，或者只希望使用一次的函数，或者可以在一行上定义的函数。让我们现在打开 IPython 并定义一个将数字平方的小函数，如pow2 = lambda x : x**2。让我们在一些数字上测试一下：

让我们回顾一下map作用于两个输入值：一个函数和给定函数可以作用的对象列表。map输出原始列表中每个元素的函数输出列表。现在让我们将我们的平方操作定义为一个匿名函数，然后将其输入到 map 中，并使用最后几个数字的列表进行检查，如map(lambda x : x**2, [2,3,4])：

我们看到map作为ElementwiseKernel！这实际上是函数式编程中的标准设计模式。现在，让我们看看reduce；它不是接收一个列表并直接输出相应列表，而是接收一个列表，在其上执行递归二进制操作，并输出一个单例。让我们通过键入reduce(lambda x, y : x + y, [1,2,3,4])来了解这种设计模式。当我们在 IPython 中键入这个时，我们将看到这将输出一个单个数字 10，这确实是1+2+3+4的和。您可以尝试用乘法替换上面的求和，并看到这确实适用于递归地将一长串数字相乘。一般来说，我们使用可结合的二进制操作进行缩减操作；这意味着，无论我们在列表的连续元素之间以何种顺序执行操作，都将始终得到相同的结果，前提是列表保持有序。（这与交换律不同。）

现在我们将看看 PyCUDA 如何处理类似于reduce的编程模式——使用并行扫描和归约内核。

并行扫描和归约内核基础

让我们看一下 PyCUDA 中一个复制 reduce 功能的基本函数——InclusiveScanKernel。（您可以在simple_scankernal0.py文件名下找到代码。）让我们执行一个在 GPU 上对一小组数字求和的基本示例：

import numpy as np
import pycuda.autoinit
from pycuda import gpuarray
from pycuda.scan import InclusiveScanKernel
seq = np.array([1,2,3,4],dtype=np.int32)
seq_gpu = gpuarray.to_gpu(seq)
sum_gpu = InclusiveScanKernel(np.int32, "a+b")
print sum_gpu(seq_gpu).get()
print np.cumsum(seq)

我们通过首先指定输入/输出类型（这里是 NumPy int32）和字符串"a+b"来构建我们的内核。在这里，InclusiveScanKernel自动在 GPU 空间中设置了名为a和b的元素，因此您可以将此字符串输入视为 Python 中的lambda a,b: a + b的类似物。我们实际上可以在这里放置任何（可结合的）二进制操作，只要我们记得用 C 语言编写它。

当我们运行sum_gpu时，我们会得到一个与输入数组大小相同的数组。数组中的每个元素表示计算中的每个步骤的值（我们可以看到，NumPy cumsum函数给出了相同的输出）。最后一个元素将是我们正在寻找的最终输出，对应于 reduce 的输出：

让我们尝试一些更具挑战性的东西；让我们找到一个float32数组中的最大值：

import numpy as np
import pycuda.autoinit
from pycuda import gpuarray
from pycuda.scan import InclusiveScanKernel
seq = np.array([1,100,-3,-10000, 4, 10000, 66, 14, 21],dtype=np.int32)
seq_gpu = gpuarray.to_gpu(seq)
max_gpu = InclusiveScanKernel(np.int32, "a > b ? a : b")
print max_gpu(seq_gpu).get()[-1]
print np.max(seq)

（您可以在名为simple_scankernal1.py的文件中找到完整的代码。）

在这里，我们所做的主要更改是用a > b ? a : b替换了a + b字符串。 （在 Python 中，这将在reduce语句中呈现为lambda a, b: max(a,b)）。在这里，我们使用了一个技巧，使用 C 语言的?运算符来给出a和b中的最大值。最后，我们显示了输出数组中结果元素的最后一个值，这将恰好是最后一个元素（我们总是可以用 Python 中的[-1]索引来检索）。

现在，让我们最后再看一个用于生成 GPU 内核的 PyCUDA 函数——ReductionKernel。实际上，ReductionKernel的作用类似于ElementwiseKernel函数，后面跟着一个并行扫描内核。哪种算法是使用ReductionKernel实现的一个好选择？首先想到的是线性代数中的点积。让我们记住计算两个向量的点积有两个步骤：

1. 将向量逐点相乘

2. 对结果的逐点乘积求和

这两个步骤也称为乘法和累加。现在让我们设置一个内核来执行这个计算：

dot_prod = ReductionKernel(np.float32, neutral="0", reduce_expr="a+b", map_expr="vec1[i]*vec2[i]", arguments="float *vec1, float *vec2") 

首先，注意我们为内核使用的数据类型（float32）。然后，我们使用arguments设置了我们的 CUDA C 内核的输入参数（这里是两个代表每个向量的浮点数组，用float *表示），并使用map_expr设置了逐点计算，这里是逐点乘法。与ElementwiseKernel一样，这是按i索引的。我们设置了reduce_expr，与InclusiveScanKernel一样。这将对数组执行元素操作的结果进行减少类型的操作。最后，我们使用neutral设置了中性元素。这是一个将作为reduce_expr的标识的元素；在这里，我们设置neutral=0，因为0在加法下始终是标识（在乘法下，1 是标识）。稍后在本书中更深入地讨论并行前缀时，我们将看到为什么我们必须设置这个。

摘要

我们首先学习了如何从 PyCUDA 查询我们的 GPU，并用此方法在 Python 中重新创建了 CUDA 的deviceQuery程序。然后我们学习了如何使用 PyCUDA 的gpuarray类及其to_gpu和get函数将 NumPy 数组传输到 GPU 的内存中。我们通过观察如何使用gpuarray对象来进行基本的 GPU 计算来感受了使用gpuarray对象的感觉，并且我们学会了使用 IPython 的prun分析器进行一些调查工作。我们发现，由于 PyCUDA 启动 NVIDIA 的nvcc编译器来编译内联 CUDA C 代码，有时在会话中首次运行 PyCUDA 的 GPU 函数时会出现一些任意的减速。然后我们学习了如何使用ElementwiseKernel函数来编译和启动逐元素操作，这些操作会自动并行化到 GPU 上。我们对 Python 中的函数式编程进行了简要回顾（特别是map和reduce函数），最后，我们介绍了如何使用InclusiveScanKernel和ReductionKernel函数在 GPU 上进行一些基本的 reduce/scan 类型计算。

现在我们已经掌握了编写和启动内核函数的绝对基础知识，我们应该意识到 PyCUDA 已经通过其模板为我们覆盖了编写内核的大部分开销。我们将在下一章学习 CUDA 内核执行的原则，以及 CUDA 如何将内核中的并发线程排列成抽象的网格和块。

问题

1. 在simple_element_kernel_example0.py中，我们在测量 GPU 计算时间时不考虑与 GPU 之间的内存传输。尝试使用 Python 时间命令测量gpuarray函数to_gpu和get的时间。考虑内存传输时间后，你会认为将这个特定函数卸载到 GPU 上值得吗？

2. 在第一章中，为什么进行 GPU 编程？，我们讨论了安德尔定律，这让我们对将程序的部分内容卸载到 GPU 上可能获得的收益有了一些了解。在本章中我们看到的两个问题，安德尔定律没有考虑到的是什么？

3. 修改gpu_mandel0.py以使用越来越小的复数格点，并将其与程序的 CPU 版本进行比较。我们可以选择足够小的格点，以至于 CPU 版本实际上比 GPU 版本更快吗？

4. 创建一个使用ReductionKernel的内核，该内核在 GPU 上获取两个相同长度的complex64数组，并返回两个数组中的绝对最大元素。

5. 如果一个gpuarray对象在 Python 中到达作用域的末尾会发生什么？

6. 你认为我们在使用ReductionKernel时为什么需要定义neutral？

7. 如果在ReductionKernel中我们设置reduce_expr ="a > b ? a : b"，并且我们正在操作 int32 类型，那么我们应该将"neutral"设置为什么？

第四章：核心、线程、块和网格

本章中，我们将看到如何编写有效的CUDA 核心。在 GPU 编程中，核心（我们可以互换使用术语CUDA 核心或核心函数）是一个可以直接从主机（CPU）启动到设备（GPU）的并行函数，而设备函数是一个只能从核心函数或另一个设备函数调用的函数。（一般来说，设备函数看起来和行为像普通的串行 C/C++函数，只是它们在 GPU 上运行，并且从核心函数并行调用。）

然后我们将了解 CUDA 如何使用线程、块和网格的概念来抽象 GPU 的一些基础技术细节（例如核心、warp 和流多处理器，我们将在本书的后面部分介绍），以及我们如何使用这些概念来减轻并行编程中的认知负担。我们将学习关于线程同步（块级和网格级），以及在 CUDA 中使用全局和共享****内存进行线程间通信。最后，我们将深入了解如何在 GPU 上实现我们自己的并行前缀类型算法（即我们在上一章中介绍的扫描/归约类型函数），这使我们能够将本章学到的所有原则付诸实践。

本章的学习成果如下：

• 理解核心和设备函数之间的区别

• 如何在 PyCUDA 中编译和启动核心，并在核心内使用设备函数

• 在启动核心的上下文中有效使用线程、块和网格，以及如何在核心内使用threadIdx和blockIdx

• 如何以及为什么在核心内同步线程，使用__syncthreads()来同步单个块中的所有线程，以及主机来同步整个块网格中的所有线程

• 如何使用设备全局和共享内存进行线程间通信

• 如何使用我们新获得的关于核心的所有知识来正确实现并行前缀和的 GPU 版本

技术要求

本章需要一台带有现代 NVIDIA GPU（2016 年以后）的 Linux 或 Windows 10 PC，并安装了所有必要的 GPU 驱动程序和 CUDA Toolkit（9.0 及以上）。还需要一个合适的 Python 2.7 安装（如 Anaconda Python 2.7），并安装了 PyCUDA 模块。

本章的代码也可以在 GitHub 上找到：

github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA

有关先决条件的更多信息，请查看本书的前言；有关软件和硬件要求，请查看github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA中的README部分。

核心

与上一章一样，我们将学习如何在 Python 代码中以内联 CUDA C 编写 CUDA 核心函数，并使用 PyCUDA 将它们启动到我们的 GPU 上。在上一章中，我们使用 PyCUDA 提供的模板来编写符合特定设计模式的核心，相比之下，我们现在将看到如何从头开始编写我们自己的核心，以便我们可以编写各种各样的核心，这些核心可能不属于 PyCUDA 涵盖的任何特定设计模式，并且我们可以更精细地控制我们的核心。当然，这些收益将以编程复杂性增加为代价；我们特别需要了解线程、块和网格及其在核心中的作用，以及如何同步我们的核心正在执行的线程，以及如何在线程之间交换数据。

让我们从简单开始，尝试重新创建我们在上一章中看到的一些逐元素操作，但这次不使用ElementwiseKernel函数；我们现在将使用SourceModule函数。这是 PyCUDA 中非常强大的函数，允许我们从头构建一个内核，所以通常最好从简单开始。

PyCUDA SourceModule 函数

我们将使用 PyCUDA 的SourceModule函数将原始内联 CUDA C 代码编译为可用的内核，我们可以从 Python 中启动。我们应该注意，SourceModule实际上将代码编译为CUDA 模块，这类似于 Python 模块或 Windows DLL，只是它包含一组编译的 CUDA 代码。这意味着我们必须使用 PyCUDA 的get_function“提取”我们想要使用的内核的引用，然后才能实际启动它。让我们从如何使用SourceModule的基本示例开始。

与以前一样，我们将从制作最简单的内核函数之一开始，即将向量乘以标量。我们将从导入开始：

import pycuda.autoinit
import pycuda.driver as drv
import numpy as np
from pycuda import gpuarray
from pycuda.compiler import SourceModule

现在我们可以立即开始编写我们的内核：

ker = SourceModule("""
__global__ void scalar_multiply_kernel(float *outvec, float scalar, float *vec)
{
 int i = threadIdx.x;
 outvec[i] = scalar*vec[i];
}
""")

因此，让我们停下来，对比一下在ElementwiseKernel中是如何完成的。首先，在 CUDA C 中声明内核函数时，我们要在前面加上__global__关键字。这将使编译器将该函数标识为内核。我们总是将其声明为void函数，因为我们总是通过传递指向一些空内存块的指针来获得输出值。我们可以像声明任何标准 C 函数一样声明参数：首先是outvec，这将是我们的输出缩放向量，当然是浮点数组指针。接下来是scalar，用一个简单的float表示；注意这不是一个指针！如果我们希望将简单的单例输入值传递给我们的内核，我们总是可以在不使用指针的情况下这样做。最后，我们有我们的输入向量vec，当然是另一个浮点数组指针。

单例输入参数可以直接从主机传递给内核函数，而无需使用指针或分配的设备内存。

让我们在继续测试之前先深入了解内核。我们记得ElementwiseKernel会自动并行化多个 GPU 线程，通过 PyCUDA 为我们设置的值i；每个单独线程的标识由threadIdx值给出，我们可以通过以下方式检索：int i = threadIdx.x;。

threadIdx用于告诉每个单独的线程其身份。这通常用于确定应在输入和输出数据数组上处理哪些值的索引。（这也可以用于使用标准 C 控制流语句（如if或switch）为特定线程分配不同的任务。）

现在，我们准备像以前一样并行执行标量乘法：outvec[i] = scalar*vec[i];。

现在，让我们测试这段代码：我们首先必须从我们刚刚使用SourceModule编译的 CUDA 模块中提取编译的内核函数的引用。我们可以使用 Python 的get_function来获取这个内核引用，如下所示：

scalar_multiply_gpu = ker.get_function("scalar_multiply_kernel")

现在，我们必须在 GPU 上放一些数据来实际测试我们的内核。让我们设置一个包含 512 个随机值的浮点数组，然后使用gpuarray.to_gpu函数将这些值复制到 GPU 的全局内存中的数组中。（我们将在 GPU 和 CPU 上将这个随机向量乘以一个标量，并查看输出是否匹配。）我们还将使用gpuarray.empty_like函数在 GPU 的全局内存中分配一块空内存块：

testvec = np.random.randn(512).astype(np.float32)
testvec_gpu = gpuarray.to_gpu(testvec)
outvec_gpu = gpuarray.empty_like(testvec_gpu)

现在，我们准备启动我们的内核。我们将标量值设置为2。（再次，由于标量是单例，我们不必将该值复制到 GPU，但是我们必须小心确保正确地进行类型转换。）在这里，我们必须使用block和grid参数明确设置线程数为512。我们现在准备启动：

scalar_multiply_gpu( outvec_gpu, np.float32(2), testvec_gpu, block=(512,1,1), grid=(1,1,1))

现在我们可以使用gpuarray输出对象中的get函数来检查输出是否与预期输出匹配，并将其与 NumPy 的allclose函数进行比较：

print "Does our kernel work correctly? : {}".format(np.allclose(outvec_gpu.get() , 2*testvec) )

（此示例的代码可在存储库中的simple_scalar_multiply_kernel.py文件中的4下找到。）

现在我们开始去掉前一章中学到的 PyCUDA 内核模板的训练轮——我们现在可以直接用纯 CUDA C 编写内核，并启动它在 GPU 上使用特定数量的线程。但是，在继续使用内核之前，我们必须更多地了解 CUDA 如何将线程结构化为抽象单位块和网格。

线程、块和网格

到目前为止，在本书中，我们一直认为线程这个术语是理所当然的。让我们退后一步，看看这究竟意味着——线程是在 GPU 的单个核心上执行的一系列指令—核心和线程不应被视为同义词！事实上，可以启动使用的线程数量比 GPU 上的核心数量多得多。这是因为，类似于英特尔芯片可能只有四个核心，但在 Linux 或 Windows 中运行数百个进程和数千个线程，操作系统的调度程序可以快速在这些任务之间切换，使它们看起来是同时运行的。GPU 以类似的方式处理线程，允许在成千上万的线程上进行无缝计算。

多个线程在 GPU 上以抽象单位块执行。您应该回忆一下我们如何从标量乘法内核中的threadIdx.x获得线程 ID；末尾有一个x，因为还有threadIdx.y和threadIdx.z。这是因为您可以在三个维度上对块进行索引，而不仅仅是一个维度。为什么我们要这样做？让我们回忆一下有关从第一章中计算 Mandelbrot 集的示例，为什么使用 GPU 编程？和第三章，使用 PyCUDA 入门。这是在二维平面上逐点计算的。因此，对于这样的算法，我们可能更倾向于在两个维度上对线程进行索引。同样，在某些情况下，使用三个维度可能是有意义的——在物理模拟中，我们可能需要在 3D 网格内计算移动粒子的位置。

块进一步以称为网格的抽象批次执行，最好将其视为块的块。与块中的线程一样，我们可以使用blockIdx.x、blockIdx.y和blockIdx.z给出的常量值在网格中的最多三个维度上对每个块进行索引。让我们看一个示例来帮助我们理解这些概念；为了简单起见，我们这里只使用两个维度。

康威的生命游戏

《生命游戏》（通常简称为 LIFE）是一种细胞自动机模拟，由英国数学家约翰·康威于 1970 年发明。听起来很复杂，但实际上非常简单——LIFE 是一个零玩家的“游戏”，由一个二维二进制格子组成，其中的“细胞”被认为是“活着的”或“死了的”。这个格子通过以下一组规则进行迭代更新：

• 任何活细胞周围少于两个活邻居的细胞会死亡

• 任何活细胞周围有两个或三个邻居的细胞会存活

• 任何活细胞周围有三个以上的邻居的细胞会死亡

• 任何死细胞周围恰好有三个邻居的细胞会复活

这四条简单的规则产生了一个复杂的模拟，具有有趣的数学特性，而且在动画时也非常美观。然而，在晶格中有大量的细胞时，它可能运行得很慢，并且通常在纯串行 Python 中编程时会导致动画不流畅。然而，这是可以并行化的，因为很明显晶格中的每个细胞可以由一个单独的 CUDA 线程管理。

我们现在将 LIFE 实现为一个 CUDA 核函数，并使用 matplotlib.animation 模块来进行动画。这对我们来说现在很有趣，因为我们将能够在这里应用我们对块和网格的新知识。

我们将首先包括适当的模块，如下所示：

import pycuda.autoinit
import pycuda.driver as drv
from pycuda import gpuarray
from pycuda.compiler import SourceModule
import numpy as np
import matplotlib.pyplot as plt 
import matplotlib.animation as animation

现在，让我们通过 SourceModule 来编写我们的核函数。我们将首先使用 C 语言的 #define 指令来设置一些我们将在整个核函数中使用的常量和宏。让我们看看我们将设置的前两个，_X 和 _Y：

ker = SourceModule("""
#define _X  ( threadIdx.x + blockIdx.x * blockDim.x )
#define _Y  ( threadIdx.y + blockIdx.y * blockDim.y )

首先让我们记住这里 #define 的工作原理——它会在编译时用定义的值（在括号中）直接替换 _X 或 _Y 的任何文本，也就是说，它为我们创建了宏。（作为个人风格的问题，我通常会在所有的 C 宏之前加上下划线。）

在 C 和 C++中，#define 用于创建宏。这意味着 #define 不会创建任何函数或设置正确的常量变量——它只允许我们在编译之前通过交换文本来在我们的代码中以简写方式编写东西。

现在，让我们具体讨论一下 _X 和 _Y 的含义——这将是单个 CUDA 线程在我们用于 LIFE 的二维晶格上的笛卡尔 x 和 y 值。我们将在一个二维网格上启动核函数，由二维块组成，这些块将对应整个细胞晶格。我们将使用线程和块常量来找到晶格上的笛卡尔点。让我们看一些图表来说明这一点。驻留在二维 CUDA 块中的线程可以被可视化如下：

此时，你可能会想知道为什么我们不在一个单独的块上启动我们的核函数，这样我们就可以将 _X 设置为 threadIdx.x，将 _Y 设置为 threadIdx.y，然后就完成了。这是由于 CUDA 对我们施加了块大小的限制——目前只支持由最多 1024 个线程组成的块。这意味着我们只能将我们的细胞晶格的尺寸最大设为 32 x 32，这将导致一个相当无聊的模拟，最好在 CPU 上完成，所以我们将在网格上启动多个块。（我们当前块的尺寸将由 blockDim.x 和 blockDim.y 给出，这将帮助我们确定目标 x 和 y 坐标，正如我们将看到的。）

同样，和之前一样，我们可以确定我们在二维网格中的块是哪个，使用 blockIdx.x 和 blockIdx.y：

在我们稍微思考一下数学之后，应该很清楚 _X 应该被定义为 (threadIdx.x + blockIdx.x * blockDim.x)，而 _Y 应该被定义为 (threadIdx.y + blockIdx.y * blockDim.y)。（添加括号是为了不干扰宏插入代码时的运算顺序。）现在，让我们继续定义剩下的宏：

#define _WIDTH  ( blockDim.x * gridDim.x )
#define _HEIGHT ( blockDim.y * gridDim.y  )

#define _XM(x)  ( (x + _WIDTH) % _WIDTH )
#define _YM(y)  ( (y + _HEIGHT) % _HEIGHT )

_WIDTH和_HEIGHT宏将分别给出我们单元格格子的宽度和高度，这应该从图表中清楚地看出。让我们讨论_XM和_YM宏。在我们的 LIFE 实现中，我们将端点“环绕”到格子的另一侧 - 例如，我们将考虑-1的x值为_WIDTH - 1，y值为-1为_HEIGHT - 1，同样地，我们将考虑_WIDTH的x值为0，y值为_HEIGHT为0。我们为什么需要这个？当我们计算给定单元格的存活邻居数时，我们可能处于某个边缘，邻居可能是外部点 - 定义这些宏来调制我们的点将自动为我们覆盖这一点。请注意，在使用 C 的模运算符之前，我们必须添加宽度或高度 - 这是因为，与 Python 不同，C 中的模运算符对于整数可以返回负值。

我们现在有一个最终的宏要定义。我们记得 PyCUDA 将二维数组作为一维指针传递到 CUDA C 中；二维数组从 Python 以按行的方式传递到一维 C 指针中。这意味着我们必须将格子上给定的笛卡尔（x，y）点转换为指向格子对应的指针中的一维点。在这里，我们可以这样做：

#define _INDEX(x,y)  ( _XM(x)  + _YM(y) * _WIDTH )

由于我们的单元格格子是按行存储的，我们必须将y值乘以宽度以偏移到对应行的点。我们现在终于可以开始我们的 LIFE 实现了。让我们从 LIFE 最重要的部分开始 - 计算给定单元格的存活邻居数。我们将使用 CUDA 设备函数来实现这一点，如下所示：

__device__ int nbrs(int x, int y, int * in)
{
     return ( in[ _INDEX(x -1, y+1) ] + in[ _INDEX(x-1, y) ] + in[ _INDEX(x-1, y-1) ] \
                   + in[ _INDEX(x, y+1)] + in[_INDEX(x, y - 1)] \
                   + in[ _INDEX(x+1, y+1) ] + in[ _INDEX(x+1, y) ] + in[ _INDEX(x+1, y-1) ] );
}

设备函数是以串行方式编写的 C 函数，由内核中的单个 CUDA 线程调用。也就是说，这个小函数将由我们内核中的多个线程并行调用。我们将把我们的单元格格子表示为 32 位整数的集合（1 表示活细胞，0 表示死细胞），所以这对我们的目的是有效的；我们只需要添加周围当前单元格的值。

CUDA 设备函数是由内核中的单个 CUDA 线程调用的串行 C 函数。虽然这些函数本身是串行的，但它们可以由多个 GPU 线程并行运行。设备函数本身不能由主机计算机启动到 GPU 上，只能由内核启动。

我们现在准备编写 LIFE 的内核实现。实际上，我们已经完成了大部分的艰苦工作 - 我们检查当前线程单元格的邻居数量，检查当前单元格是生还是死，然后使用适当的 switch-case 语句来根据 LIFE 的规则确定下一次迭代的状态。我们将使用两个整数指针数组作为内核 - 一个将用作输入参考上一次迭代（lattice），另一个将用作输出我们将计算的迭代（lattice_out）的参考。

__global__ void conway_ker(int * lattice_out, int * lattice  )
{
   // x, y are the appropriate values for the cell covered by this thread
   int x = _X, y = _Y;

   // count the number of neighbors around the current cell
   int n = nbrs(x, y, lattice);

    // if the current cell is alive, then determine if it lives or dies for the next generation.
    if ( lattice[_INDEX(x,y)] == 1)
       switch(n)
       {
          // if the cell is alive: it remains alive only if it has 2 or 3 neighbors.
          case 2:
          case 3: lattice_out[_INDEX(x,y)] = 1;
                  break;
          default: lattice_out[_INDEX(x,y)] = 0;                   
       }
    else if( lattice[_INDEX(x,y)] == 0 )
         switch(n)
         {
            // a dead cell comes to life only if it has 3 neighbors that are alive.
            case 3: lattice_out[_INDEX(x,y)] = 1;
                    break;
            default: lattice_out[_INDEX(x,y)] = 0;         
         }

}
""")

conway_ker = ker.get_function("conway_ker")

我们记得用三重括号关闭内联 CUDA C 段落，然后用get_function获取对我们的 CUDA C 内核的引用。由于内核只会一次更新格子，我们将在 Python 中设置一个简短的函数，它将涵盖更新格子的所有开销以用于动画：

def update_gpu(frameNum, img, newLattice_gpu, lattice_gpu, N):    

frameNum参数只是 Matplotlib 动画模块对于我们可以忽略的更新函数所需的一个值，而img将是我们单元格格子的代表图像，这是动画模块所需的，将被迭代显示。

让我们关注另外三个参数—newLattice_gpu和lattice_gpu将是我们保持持久的 PyCUDA 数组，因为我们希望避免在 GPU 上重新分配内存块。lattice_gpu将是细胞数组的当前一代，对应于内核中的lattice参数，而newLattice_gpu将是下一代晶格。N将指示晶格的高度和宽度（换句话说，我们将使用N x N晶格）。

我们使用适当的参数启动内核，并设置块和网格大小如下：

    conway_ker(newLattice_gpu, lattice_gpu, grid=(N/32,N/32,1), block=(32,32,1) )    

我们将块大小设置为 32 x 32，使用(32, 32, 1)；因为我们只使用两个维度来表示我们的细胞晶格，所以我们可以将z维度设置为 1。请记住，块的线程数限制为 1,024 个线程—32 x 32 = 1024，所以这样可以工作。（请记住，32 x 32 没有什么特别之处；如果需要，我们可以使用 16 x 64 或 10 x 10 等值，只要总线程数不超过 1,024。）

CUDA 块中的线程数最多为 1,024。

现在我们来看一下网格值—在这里，因为我们使用 32 的维度，很明显N（在这种情况下）应该是 32 的倍数。这意味着在这种情况下，我们只能使用 64 x 64、96 x 96、128 x 128 和 1024 x 1024 等晶格。同样，如果我们想使用不同大小的晶格，那么我们将不得不改变块的维度。（如果这不太清楚，请查看之前的图表，并回顾一下我们如何在内核中定义宽度和高度宏。）

现在我们可以使用get()函数从 GPU 的内存中获取最新生成的晶格，并为我们的动画设置图像数据。最后，我们使用 PyCUDA 切片操作[:]将新的晶格数据复制到当前数据中，这将复制 GPU 上先前分配的内存，这样我们就不必重新分配了：

    img.set_data(newLattice_gpu.get() )    
    lattice_gpu[:] = newLattice_gpu[:]

    return img

让我们设置一个大小为 256 x 256 的晶格。现在我们将使用numpy.random模块中的 choice 函数为我们的晶格设置初始状态。我们将随机用 1 和 0 填充一个N x N的整数图表；通常，如果大约 25%的点是 1，其余的是 0，我们可以生成一些有趣的晶格动画，所以我们就这样做吧：

if __name__ == '__main__':
    # set lattice size
    N = 256

    lattice = np.int32( np.random.choice([1,0], N*N, p=[0.25, 0.75]).reshape(N, N) )
    lattice_gpu = gpuarray.to_gpu(lattice)

最后，我们可以使用适当的gpuarray函数在 GPU 上设置晶格，并相应地设置 Matplotlib 动画，如下所示：

lattice_gpu = gpuarray.to_gpu(lattice)
    lattice_gpu = gpuarray.to_gpu(lattice)
    newLattice_gpu = gpuarray.empty_like(lattice_gpu) 

    fig, ax = plt.subplots()
    img = ax.imshow(lattice_gpu.get(), interpolation='nearest')
    ani = animation.FuncAnimation(fig, update_gpu, fargs=(img, newLattice_gpu, lattice_gpu, N, ) , interval=0, frames=1000, save_count=1000) 

    plt.show()

现在我们可以运行我们的程序并享受展示（代码也可以在 GitHub 存储库的4目录下的conway_gpu.py文件中找到）：

线程同步和互通

现在我们将讨论 GPU 编程中的两个重要概念—线程同步和线程互通。有时，我们需要确保每个线程在继续任何进一步的计算之前都已经到达了代码中完全相同的行；我们称之为线程同步。同步与线程互通相辅相成，也就是说，不同的线程之间传递和读取输入；在这种情况下，我们通常希望确保所有线程在传递数据之前都处于计算的相同步骤。我们将从学习 CUDA __syncthreads设备函数开始，该函数用于同步内核中的单个块。

使用 __syncthreads()设备函数

在我们之前的康威生命游戏的示例中，我们的内核每次由主机启动时只更新了晶格一次。在这种情况下，同步所有在启动的内核中的线程没有问题，因为我们只需要处理已经准备好的晶格的上一个迭代。

现在假设我们想做一些稍微不同的事情——我们想重新编写我们的内核，以便在给定的细胞点阵上执行一定数量的迭代，而不是由主机一遍又一遍地重新启动。这一开始可能看起来很琐碎——一个天真的解决方案将是只需在内联conway_ker内核中放置一个整数参数来指示迭代次数和一个for循环，进行一些额外的琐碎更改，然后就完成了。

然而，这引发了竞争条件的问题；这是多个线程读取和写入相同内存地址以及由此可能产生的问题。我们的旧conway_ker内核通过使用两个内存数组来避免这个问题，一个严格用于读取，一个严格用于每次迭代写入。此外，由于内核只执行单次迭代，我们实际上是在使用主机来同步线程。

我们希望在 GPU 上进行多次完全同步的 LIFE 迭代；我们还希望使用单个内存数组来存储点阵。我们可以通过使用 CUDA 设备函数__syncthreads()来避免竞争条件。这个函数是一个块级同步屏障——这意味着在一个块内执行的每个线程在到达__syncthreads()实例时都会停止，并等待直到同一块内的每个其他线程都到达__syncthreads()的同一调用，然后线程才会继续执行后续的代码行。

__syncthreads()只能同步单个 CUDA 块内的线程，而不能同步 CUDA 网格内的所有线程！

让我们现在创建我们的新内核；这将是之前 LIFE 内核的修改，它将执行一定数量的迭代然后停止。这意味着我们不会将其表示为动画，而是作为静态图像，因此我们将在开始时加载适当的 Python 模块。（此代码也可在 GitHub 存储库的conway_gpu_syncthreads.py文件中找到）：

import pycuda.autoinit
import pycuda.driver as drv
from pycuda import gpuarray
from pycuda.compiler import SourceModule
import numpy as np
import matplotlib.pyplot as plt 

现在，让我们再次设置计算 LIFE 的内核：

ker = SourceModule("""

当然，我们的 CUDA C 代码将放在这里，这将大致与以前相同。我们只需要对内核进行一些更改。当然，我们可以保留设备函数nbrs。在我们的声明中，我们将只使用一个数组来表示细胞点阵。我们可以这样做，因为我们将使用适当的线程同步。我们还必须用一个整数表示迭代次数。我们设置参数如下：

__global__ void conway_ker(int * lattice, int iters) {

我们将继续与以前类似，只是使用for循环进行迭代：

 int x = _X, y = _Y; 
 for (int i = 0; i < iters; i++)
 {
     int n = nbrs(x, y, lattice); 
     int cell_value;

让我们回想一下以前，我们直接在数组中设置新的细胞点阵值。在这里，我们将在cell_value变量中保存值，直到块内的所有线程都同步。我们以前也是类似地进行，使用__syncthreads阻止执行，直到确定了当前迭代的所有新细胞值，然后才在点阵数组中设置值：

 if ( lattice[_INDEX(x,y)] == 1)
 switch(n)
 {
 // if the cell is alive: it remains alive only if it has 2 or 3 neighbors.
 case 2:
 case 3: cell_value = 1;
 break;
 default: cell_value = 0; 
 }
 else if( lattice[_INDEX(x,y)] == 0 )
 switch(n)
 {
 // a dead cell comes to life only if it has 3 neighbors that are alive.
 case 3: cell_value = 1;
 break;
 default: cell_value = 0; 
 } 
 __syncthreads();
 lattice[_INDEX(x,y)] = cell_value; 
 __syncthreads();
 } 
}
""")

我们现在将像以前一样启动内核并显示输出，迭代点阵 100 万次。请注意，由于每个块的线程限制为 1,024 个，我们在网格中只使用一个块，大小为 32 x 32。（再次强调，__syncthreads仅在块中的所有线程上工作，而不是在网格中的所有线程上工作，这就是为什么我们在这里限制自己使用单个块的原因）：

conway_ker = ker.get_function("conway_ker")
if __name__ == '__main__':
 # set lattice size
 N = 32
 lattice = np.int32( np.random.choice([1,0], N*N, p=[0.25, 0.75]).reshape(N, N) )
 lattice_gpu = gpuarray.to_gpu(lattice)
 conway_ker(lattice_gpu, np.int32(1000000), grid=(1,1,1), block=(32,32,1))
 fig = plt.figure(1)
 plt.imshow(lattice_gpu.get())

当我们运行程序时，我们将得到以下所需的输出（这是随机 LIFE 点阵在一百万次迭代后会收敛到的结果！）：

使用共享内存

从先前的例子中，我们可以看到内核中的线程可以使用 GPU 全局内存中的数组进行相互通信；虽然可以使用全局内存进行大多数操作，但使用共享内存可以加快速度。这是一种专门用于单个 CUDA 块内线程相互通信的内存类型；与全局内存相比，使用共享内存的优势在于纯线程间通信速度更快。不过，与全局内存相反，存储在共享内存中的内存不能直接被主机访问——共享内存必须首先由内核自己复制回全局内存。

在继续之前，让我们先退一步思考一下我们的意思。让我们看看我们刚刚看到的迭代 LIFE 内核中声明的一些变量。首先看看x和y，这两个整数保存着特定线程单元的笛卡尔坐标。请记住，我们正在使用_X和_Y宏设置它们的值。（尽管编译器优化，我们希望将这些值存储在变量中以减少计算，因为直接使用_X和_Y将在我们的代码中引用这些宏时每次重新计算x和y的值）：

 int x = _X, y = _Y; 

我们注意到，对于每个单个线程，点阵中将对应于x和y的唯一笛卡尔点。同样，我们使用一个变量n，它声明为int n = nbrs(x, y, lattice);，来表示特定单元周围的存活邻居的数量。这是因为，当我们通常在 CUDA 中声明变量时，它们默认是每个单独线程的本地变量。请注意，即使我们在线程内部声明数组如int a[10];，也将有一个大小为 10 的数组，它是每个线程的本地数组。

本地线程数组（例如，在内核内部声明int a[10];）和指向全局 GPU 内存的指针（例如，以int * b形式作为内核参数传递的值）可能看起来和行为类似，但实际上非常不同。对于内核中的每个线程，将有一个单独的a数组，其他线程无法读取，但将有一个单独的b，它将保存相同的值，并且对所有线程都是同样可访问的。

我们准备使用共享内存。这使我们能够声明在单个 CUDA 块内的线程之间共享的变量和数组。这种内存比使用全局内存指针（我们到目前为止一直在使用的）要快得多，同时减少了指针分配内存的开销。

假设我们想要一个大小为 10 的共享整数数组。我们声明如下——__shared__ int a[10] 。请注意，我们不必局限于数组；我们可以按如下方式创建共享的单例变量：__shared__ int x。

让我们重新编写 LIFE 的迭代版本的一些行，以利用共享内存。首先，让我们将输入指针重命名为p_lattice，这样我们可以在我们的共享数组上使用这个变量名，并在我们的代码中懒惰地保留所有对lattice的引用。由于我们将坚持使用 32 x 32 个单元的点阵，我们设置新的共享lattice数组如下：

__global__ void conway_ker_shared(int * p_lattice, int iters)
{
 int x = _X, y = _Y;
 __shared__ int lattice[32*32];

现在我们必须将全局内存p_lattice数组中的所有值复制到lattice中。我们将以完全相同的方式索引我们的共享数组，因此我们可以在这里使用我们旧的_INDEX宏。请注意，在复制后我们确保在我们继续 LIFE 算法之前放置__syncthreads()，以确保所有对 lattice 的内存访问完全完成：

 lattice[_INDEX(x,y)] = p_lattice[_INDEX(x,y)];
 __syncthreads();

内核的其余部分与以前完全相同，只是我们必须将共享的点阵复制回 GPU 数组。我们这样做，然后关闭内联代码：

 __syncthreads();
 p_lattice[_INDEX(x,y)] = lattice[_INDEX(x,y)];
 __syncthreads();
} """)

现在我们可以像以前一样运行，使用完全相同的测试代码。（此示例可以在 GitHub 存储库中的conway_gpu_syncthreads_shared.py中找到。）

并行前缀算法

现在，我们将利用我们对 CUDA 核心的新知识来实现并行前缀算法，也称为扫描设计模式。我们已经在上一章中以 PyCUDA 的InclusiveScanKernel和ReductionKernel函数的形式看到了这种简单的例子。现在让我们更详细地了解这个想法。

这个设计模式的核心动机是，我们有一个二元运算符，也就是说，一个作用于两个输入值并给出一个输出值的函数（比如—+，，（最大值），（最小值）），和元素的集合，我们希望从中高效地计算。此外，我们假设我们的二元运算符是可结合的—这意味着，对于任意三个元素x、y和z，我们总是有:。

我们希望保留部分结果，也就是n-1个子计算—。并行前缀算法的目的是高效地产生这些n个和。在串行操作中，通常需要O(n)的时间来产生这些n个和，我们希望降低时间复杂度。

当使用术语“并行前缀”或“扫描”时，通常意味着一个产生所有这些n个结果的算法，而“减少”/“归约”通常意味着只产生单个最终结果，。（这是 PyCUDA 的情况。）

实际上，并行前缀算法有几种变体，我们将首先从最简单（也是最古老的）版本开始，这就是所谓的天真并行前缀算法。

天真并行前缀算法

这个算法是原始版本的天真并行前缀算法；这个算法是“天真的”，因为它假设给定n个输入元素，，进一步假设n是二进制的（也就是说，对于某个正整数k），我们可以在n个处理器（或n个线程）上并行运行算法。显然，这将对我们可以处理的集合的基数n施加严格的限制。然而，只要满足这些条件，我们就有一个很好的结果，即其计算时间复杂度仅为O(log n)。我们可以从算法的伪代码中看到这一点。在这里，我们将用表示输入值，用表示输出值：

input: x0, ..., xn-1 initialize:
for k=0 to n-1:
    yk := xk begin:
parfor i=0 to n-1 :
    for j=0 to log2(n):
        if i >= 2j :
            yi := yi  yi - 2^j else:
            continue
        end if
    end for
end parfor
end
output: y0, ..., yn-1

现在，我们可以清楚地看到这将花费O(log n)的渐近时间，因为外部循环是在parfor上并行化的，内部循环需要log2。经过几分钟的思考，很容易看出y[i]值将产生我们期望的输出。

现在让我们开始我们的实现；在这里，我们的二元运算符将简单地是加法。由于这个例子是说明性的，这个核心代码将严格地针对 1,024 个线程。

让我们先设置好头文件，然后开始编写我们的核心代码：

import pycuda.autoinit
import pycuda.driver as drv
import numpy as np
from pycuda import gpuarray
from pycuda.compiler import SourceModule
from time import time

naive_ker = SourceModule("""
__global__ void naive_prefix(double *vec, double *out)
{
     __shared__ double sum_buf[1024]; 
     int tid = threadIdx.x; 
     sum_buf[tid] = vec[tid];

因此，让我们看看我们有什么：我们将我们的输入元素表示为双精度 GPU 数组，即double *vec，并用double *out表示输出值。我们声明一个共享内存sum_buf数组，我们将用它来计算我们的输出。现在，让我们看看算法本身的实现：

 int iter = 1;
 for (int i=0; i < 10; i++)
 {
     __syncthreads();
     if (tid >= iter )
     {
         sum_buf[tid] = sum_buf[tid] + sum_buf[tid - iter]; 
     } 
     iter *= 2;
 }
 __syncthreads();

当然，这里没有parfor，它是隐式的，通过tid变量表示线程编号。我们还可以通过从初始化为 1 的变量开始，然后在每次 i 的迭代中乘以 2 来省略log[2]和2^i的使用。（请注意，如果我们想更加技术化，我们可以使用位移运算符来实现这一点。）我们将i的迭代限制在 10 次，因为2¹⁰ = 1024。现在我们将结束我们的新内核如下：

 __syncthreads();
 out[tid] = sum_buf[tid];
 __syncthreads();

}
""")
naive_gpu = naive_ker.get_function("naive_prefix")

现在让我们看一下内核后面的测试代码：

if __name__ == '__main__':
 testvec = np.random.randn(1024).astype(np.float64)
 testvec_gpu = gpuarray.to_gpu(testvec)

 outvec_gpu = gpuarray.empty_like(testvec_gpu)
 naive_gpu( testvec_gpu , outvec_gpu, block=(1024,1,1), grid=(1,1,1))

 total_sum = sum( testvec)
 total_sum_gpu = outvec_gpu[-1].get()

 print "Does our kernel work correctly? : {}".format(np.allclose(total_sum_gpu , total_sum) )

我们只关心输出中的最终总和，我们使用outvec_gpu[-1].get()来检索它，需要回想一下，"-1"索引给出了 Python 数组中的最后一个成员。这将是vec中每个元素的总和；部分总和在outvec_gpu的先前值中。（此示例可以在 GitHub 存储库中的naive_prefix.py文件中看到。）

由于其性质，并行前缀算法必须在n个线程上运行，对应于一个大小为 n 的数组，其中n是二进制的（这意味着n是 2 的某个幂）。然而，我们可以将这个算法扩展到任意非二进制大小，假设我们的运算符具有单位元素（或等效地，中性元素）——也就是说，存在某个值e，使得对于任何x值，我们有 。在运算符为+的情况下，单位元素是 0；在运算符为  的情况下，它是 1；然后我们只需用一系列e值填充  的元素，以便我们有新集合的二进制基数 。

包含与排他前缀

让我们停下来，做一个非常微妙但非常重要的区分。到目前为止，我们一直关心接受形式为  的输入，并产生形式为  的总和数组作为输出。产生这种输出的前缀算法称为包含；在包含前缀算法的情况下，每个索引处的相应元素包含在输出数组的相同索引处的总和中。这与排他前缀算法形成对比。排他前缀算法不同之处在于，它同样接受形式为  的n个输入值，并产生长度为n的输出数组 。

这很重要，因为一些高效的前缀算法的变体天生就是排他的。我们将在下一个小节中看到一个例子。

请注意，排他算法产生的输出与包含算法几乎相同，只是右移并省略了最后一个值。因此，我们可以从任一算法中轻松获得等效输出，只要我们保留  的副本。

一个高效的并行前缀算法

在我们继续使用新算法之前，我们将从两个角度看待朴素算法。在理想情况下，计算时间复杂度为O(log n)，但这仅当我们有足够数量的处理器用于我们的数据集时成立；当数据集的基数（元素数量）n远大于处理器数量时，这将成为O(n log n)时间复杂度的算法。

让我们定义一个与我们的二进制运算符 相关的新概念——这里并行算法的工作是执行期间所有线程对此运算符的调用次数。同样，跨度是线程在内核执行期间进行调用的次数；而整个算法的跨度与每个单独线程的最长跨度相同，这将告诉我们总执行时间。

我们寻求特别减少算法在所有线程中执行的工作量，而不仅仅是关注跨度。在简单前缀的情况下，所需的额外工作在可用处理器数量不足时会花费更多时间；这些额外工作将溢出到有限数量的可用处理器中。

我们将介绍一种新的算法，它是工作高效的，因此更适合有限数量的处理器。这包括两个独立的部分——向上扫描（或减少）阶段和向下扫描阶段。我们还应该注意，我们将看到的算法是一种独占前缀算法。

向上扫描阶段类似于单个减少操作，以产生由减少算法给出的值，即  ；在这种情况下，我们保留所需的部分和（）以实现最终结果。然后，向下扫描阶段将对这些部分和进行操作，并给出最终结果。让我们看一些伪代码，从向上扫描阶段开始。（接下来的小节将立即深入到伪代码的实现中。）

工作高效的并行前缀（向上扫描阶段）

这是向上扫描的伪代码。（注意parfor覆盖j变量，这意味着此代码块可以并行化，由j索引的线程）：

input: x0, ..., xn-1initialize:
    for i = 0 to n - 1:
        yi := xi
begin:
for k=0 to log2(n) - 1:
    parfor j=0 to n - 1: 
        if j is divisible by 2k+1:
            yj+2^(k+1)-1 = yj+2^k-1  yj +2^(k+1) -1else:
            continue
end
output: y0, ..., yn-1

工作高效的并行前缀（向下扫描阶段）

现在让我们继续向下扫描，它将操作向上扫描的输出：

input: x0, ..., xn-1 initialize:
    for i = 0 to n - 2:
        yi := xi
    yn-1 := 0
begin:
for k = log2(n) - 1 to 0:
    parfor j = 0 to n - 1: 
        if j is divisible by 2k+1:
            temp := yj+2^k-1
            yj+2^k-1 := yj+2^(k+1)-1
            yj+2^(k+1)-1 := yj+2^(k+1)-1  temp
        else:
            continue
end
output: y0 , y1 , ..., yn-1

工作高效的并行前缀 — 实现

作为本章的压轴，我们将编写该算法的实现，该算法可以在大于 1,024 的任意大小的数组上操作。这意味着这将在网格和块上操作；因此，我们将不得不使用主机进行同步；此外，这将要求我们为向上扫描和向下扫描阶段实现两个单独的内核，这将作为两个阶段的parfor循环，以及作为向上和向下扫描的外部for循环的 Python 函数。

让我们从向上扫描内核开始。由于我们将从主机迭代重新启动此内核，我们还需要一个指示当前迭代（k）的参数。我们将使用两个数组进行计算，以避免竞争条件——x（用于当前迭代）和x_old（用于之前的迭代）。我们声明内核如下：

up_ker = SourceModule("""
__global__ void up_ker(double *x, double *x_old, int k)
{

现在让我们设置tid变量，它将是网格中所有块中所有线程中当前线程的标识。我们使用与我们之前看到的康威生命游戏的原始网格级实现相同的技巧：

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

我们现在将使用 C 位移运算符直接从k生成 2^k 和 2^(k+1 )。我们现在将j设置为tid乘以_2k1—这将使我们能够删除“如果j可被 2^(k+1)整除”的部分，如伪代码中所示，从而使我们只启动所需数量的线程：

 int _2k = 1 << k;
 int _2k1 = 1 << (k+1);

 int j = tid* _2k1;

我们可以使用 CUDA C 中的左位移运算符（<<）轻松生成二进制（2 的幂）整数。请记住，整数 1（即 2⁰）表示为 0001，2（2¹）表示为 0010，4（2²）表示为 0100，依此类推。因此，我们可以使用1 << k操作计算 2^k。

现在我们可以运行上扫描阶段的单行代码，注意j确实可以被 2^(k+1)整除，因为它的构造方式是这样的：

 x[j + _2k1 - 1] = x_old[j + _2k -1 ] + x_old[j + _2k1 - 1];
}
""")

我们已经编写完我们的核心代码了！但这当然不是完整的上扫描实现。我们还需要在 Python 中完成其余部分。让我们拿到我们的核心代码并开始实现。这基本上是按照伪代码进行的，我们应该记住，我们通过使用[:]从x_gpu复制到x_old_gpu来更新x_old_gpu，这将保留内存分配，并仅复制新数据而不是重新分配。还要注意，我们根据要启动的线程数量设置我们的块和网格大小 - 我们尝试保持我们的块大小为 32 的倍数（这是本文中的经验法则，我们在第十一章中详细介绍为什么我们特别使用 32，CUDA 性能优化）。我们应该在文件开头加上from __future__ import division，因为我们将使用 Python 3 风格的除法来计算我们的块和核心大小。

有一点需要提到的是，我们假设x的长度是二进制长度 32 或更大 - 如果您希望将其操作在其他大小的数组上，可以通过用零填充我们的数组来轻松修改，然而：

up_gpu = up_ker.get_function("up_ker")

def up_sweep(x):
    x = np.float64(x)
    x_gpu = gpuarray.to_gpu(np.float64(x) )
    x_old_gpu = x_gpu.copy()
    for k in range( int(np.log2(x.size) ) ) : 
        num_threads = int(np.ceil( x.size / 2**(k+1)))
        grid_size = int(np.ceil(num_threads / 32))

        if grid_size > 1:
            block_size = 32
        else:
            block_size = num_threads

        up_gpu(x_gpu, x_old_gpu, np.int32(k) , block=(block_size,1,1), grid=(grid_size,1,1))
        x_old_gpu[:] = x_gpu[:]

    x_out = x_gpu.get()
    return(x_out)

现在我们将开始编写下扫描。同样，让我们从核心开始，它将具有伪代码中内部parfor循环的功能。它与之前类似 - 再次使用两个数组，因此在这里使用temp变量与伪代码中是不必要的，我们再次使用位移运算符来获得 2^k 和 2^(k+1)的值。我们计算j与之前类似：

down_ker = SourceModule("""
__global__ void down_ker(double *y, double *y_old, int k)
{
 int j = blockIdx.x*blockDim.x + threadIdx.x;

 int _2k = 1 << k;
 int _2k1 = 1 << (k+1);

 int j = tid*_2k1;

 y[j + _2k - 1 ] = y_old[j + _2k1 - 1];
 y[j + _2k1 - 1] = y_old[j + _2k1 - 1] + y_old[j + _2k - 1];
}
""")

down_gpu = down_ker.get_function("down_ker")

现在我们可以编写一个 Python 函数，该函数将迭代地启动核心，这对应于下扫描阶段的外部for循环。这类似于上扫描阶段的 Python 函数。从伪代码中看到的一个重要区别是，我们必须从外部for循环中的最大值迭代到最小值；我们可以使用 Python 的reversed函数来做到这一点。现在我们可以实现下扫描阶段：

def down_sweep(y):
    y = np.float64(y)
    y[-1] = 0
    y_gpu = gpuarray.to_gpu(y)
    y_old_gpu = y_gpu.copy()
    for k in reversed(range(int(np.log2(y.size)))):
        num_threads = int(np.ceil( y.size / 2**(k+1)))
        grid_size = int(np.ceil(num_threads / 32))

        if grid_size > 1:
            block_size = 32
        else:
            block_size = num_threads

        down_gpu(y_gpu, y_old_gpu, np.int32(k), block=(block_size,1,1), grid=(grid_size,1,1))
        y_old_gpu[:] = y_gpu[:]
    y_out = y_gpu.get()
    return(y_out)

在实现了上扫描和下扫描阶段之后，我们的最后任务是轻而易举地完成：

def efficient_prefix(x):
        return(down_sweep(up_sweep(x)))

我们现在已经完全实现了一个与主机同步的高效并行前缀算法！（这个实现可以在存储库中的work-efficient_prefix.py文件中找到，还有一些测试代码。）

总结

我们从康威的生命游戏实现开始，这给了我们一个关于 CUDA 核心的许多线程是如何在块-网格张量类型结构中组织的想法。然后，我们通过 CUDA 函数__syncthreads()深入研究了块级同步，以及通过使用共享内存进行块级线程互通；我们还看到单个块有一定数量的线程，我们可以操作，所以在创建将使用多个块跨越更大网格的核心时，我们必须小心使用这些功能。

我们概述了并行前缀算法的理论，并最后实现了一个天真的并行前缀算法，作为一个单个核心，可以操作大小受限的数组，该数组与___syncthreads同步，并在内部执行for和parfor循环，并且实现了一个高效的并行前缀算法，该算法跨两个核心和三个 Python 函数实现，可以操作任意大小的数组，核心充当算法的内部parfor循环，而 Python 函数有效地充当外部for循环并同步核心启动。

问题

1. 更改simple_scalar_multiply_kernel.py中的随机向量，使其长度为 10,000，并修改内核定义中的i索引，以便它可以在网格形式的多个块中使用。看看现在是否可以通过将块和网格参数设置为block=(100,1,1)和grid=(100,1,1)来启动这个内核超过 10,000 个线程。

2. 在上一个问题中，我们启动了一个同时使用 10,000 个线程的内核；截至 2018 年，没有 NVIDIA GPU 具有超过 5,000 个核心。为什么这仍然有效并产生预期的结果？

3. 天真的并行前缀算法在数据集大小为n的情况下，时间复杂度为 O(log n)，假设我们有n个或更多处理器。假设我们在具有 640 个核心的 GTX 1050 GPU 上使用天真的并行前缀算法。在n >> 640的情况下，渐近时间复杂度会变成什么？

4. 修改naive_prefix.py以在大小为 1,024 的数组上运行（可能是非二进制的）。

5. __syncthreads() CUDA 设备函数只在单个块中同步线程。我们如何在整个网格的所有块中同步所有线程？

6. 您可以通过这个练习说服自己，第二个前缀和算法确实比天真的前缀和算法更有效。假设我们有一个大小为 32 的数据集。在这种情况下，第一个和第二个算法所需的“加法”操作的确切数量是多少？

7. 在高效并行前缀的实现中，我们使用 Python 函数来迭代我们的内核并同步结果。为什么我们不能在内核中使用for循环，同时小心地使用__syncthreads()？

8. 为什么在 CUDA C 中实现天真的并行前缀在一个单独的内核中处理自己的同步比使用两个内核和 Python 函数实现高效并行前缀更有意义，并且让主机处理同步更有意义？

第五章：流、事件、上下文和并发性

在之前的章节中，我们看到在与 GPU 交互时，主机执行的两个主要操作是：

• 将内存数据从 GPU 复制到 GPU，以及从 GPU 复制到内存

• 启动内核函数

我们知道，在单个内核中，其许多线程之间存在一定程度的并发性；然而，在多个内核和 GPU 内存操作之间，还存在另一种并发性。这意味着我们可以同时启动多个内存和内核操作，而无需等待每个操作完成。然而，另一方面，我们必须有一定的组织能力，以确保所有相互依赖的操作都得到同步；这意味着我们不应该在其输入数据完全复制到设备内存之前启动特定的内核，或者在内核执行完成之前，不应该将已启动内核的输出数据复制到主机。

为此，我们有所谓的CUDA 流—流是在 GPU 上按顺序运行的一系列操作。单独一个流是没有用的—重点是通过使用多个流在主机上发出的 GPU 操作来获得并发性。这意味着我们应该交错启动与不同流对应的 GPU 操作，以利用这个概念。

我们将在本章中广泛涵盖流的概念。此外，我们还将研究事件，这是流的一个特性，用于精确计时内核，并指示主机在给定流中已完成哪些操作。

最后，我们将简要介绍 CUDA 上下文。上下文可以被视为类似于操作系统中的进程，因为 GPU 将每个上下文的数据和内核代码隔离并封装起来，使其与当前存在于 GPU 上的其他上下文相分离。我们将在本章末尾看到这方面的基础知识。

本章的学习成果如下：

• 理解设备和流同步的概念

• 学习如何有效地使用流来组织并发的 GPU 操作

• 学习如何有效地使用 CUDA 事件

• 理解 CUDA 上下文

• 学习如何在给定上下文中显式同步

• 学习如何显式创建和销毁 CUDA 上下文

• 学习如何使用上下文允许主机上的多个进程和线程共享 GPU 使用

技术要求

本章需要一台带有现代 NVIDIA GPU（2016 年以后）的 Linux 或 Windows 10 PC，并安装了所有必要的 GPU 驱动程序和 CUDA Toolkit（9.0 及以上）。还需要一个合适的 Python 2.7 安装（如 Anaconda Python 2.7），并安装了 PyCUDA 模块。

本章的代码也可以在 GitHub 上找到：

github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA

有关先决条件的更多信息，请查看本书的前言，有关软件和硬件要求，请查看github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA中的 README。

CUDA 设备同步

在我们可以使用 CUDA 流之前，我们需要了解设备同步的概念。这是一种操作，其中主机阻塞任何进一步的执行，直到发给 GPU 的所有操作（内存传输和内核执行）都已完成。这是为了确保依赖于先前操作的操作不会被无序执行—例如，确保 CUDA 内核启动在主机尝试读取其输出之前已完成。

在 CUDA C 中，设备同步是通过cudaDeviceSynchronize函数执行的。这个函数有效地阻止了主机上的进一步执行，直到所有 GPU 操作完成。cudaDeviceSynchronize是如此基本，以至于它通常是 CUDA C 大多数书籍中最早涉及的主题之一-我们还没有看到这一点，因为 PyCUDA 已经在需要时自动为我们调用了这个函数。让我们看一个 CUDA C 代码的例子，看看如何手动完成这个操作：

// Copy an array of floats from the host to the device.
cudaMemcpy(device_array, host_array, size_of_array*sizeof(float), cudaMemcpyHostToDevice);
// Block execution until memory transfer to device is complete.
cudaDeviceSynchronize();
// Launch CUDA kernel.
Some_CUDA_Kernel <<< block_size, grid_size >>> (device_array, size_of_array);
// Block execution until GPU kernel function returns.
cudaDeviceSynchronize();
// Copy output of kernel to host.
cudaMemcpy(host_array,  device_array, size_of_array*sizeof(float), cudaMemcpyDeviceToHost); // Block execution until memory transfer to host is complete.
cudaDeviceSynchronize();

在这段代码块中，我们看到我们必须在每个单独的 GPU 操作之后直接与设备进行同步。如果我们只需要一次调用单个 CUDA 内核，就像这里看到的那样，这是可以的。但是，如果我们想要同时启动多个独立的内核和操作不同数据数组的内存操作，跨整个设备进行同步将是低效的。在这种情况下，我们应该跨多个流进行同步。我们现在就来看看如何做到这一点。

使用 PyCUDA 流类

我们将从一个简单的 PyCUDA 程序开始；它只是生成一系列随机的 GPU 数组，对每个数组进行简单的内核处理，然后将数组复制回主机。然后我们将修改它以使用流。请记住，这个程序将完全没有任何意义，只是为了说明如何使用流以及一些基本的性能提升。（这个程序可以在 GitHub 存储库的5目录下的multi-kernel.py文件中看到。）

当然，我们将首先导入适当的 Python 模块，以及time函数：

import pycuda.autoinit
import pycuda.driver as drv
from pycuda import gpuarray
from pycuda.compiler import SourceModule
import numpy as np
from time import time

我们现在将指定要处理多少个数组-在这里，每个数组将由不同的内核启动进行处理。我们还指定要生成的随机数组的长度，如下所示：

num_arrays = 200
array_len = 1024**2

我们现在有一个在每个数组上操作的内核；它只是迭代数组中的每个点，并将其乘以 2 再除以 2，重复 50 次，最终保持数组不变。我们希望限制每个内核启动将使用的线程数，这将帮助我们在 GPU 上的许多内核启动之间获得并发性，以便我们可以让每个线程通过for循环迭代数组的不同部分。（再次提醒，这个内核函数除了用来学习流和同步之外，完全没有任何用处！）如果每个内核启动使用的线程太多，以后获得并发性将更加困难：

ker = SourceModule(""" 
__global__ void mult_ker(float * array, int array_len)
{
     int thd = blockIdx.x*blockDim.x + threadIdx.x;
     int num_iters = array_len / blockDim.x;

     for(int j=0; j < num_iters; j++)
     {
         int i = j * blockDim.x + thd;

         for(int k = 0; k < 50; k++)
         {
              array[i] *= 2.0;
              array[i] /= 2.0;
         }
     }
}
""")

mult_ker = ker.get_function('mult_ker')

现在，我们将生成一些随机数据数组，将这些数组复制到 GPU，迭代地在每个数组上启动我们的内核，然后将输出数据复制回主机，并使用 NumPy 的allclose函数来断言它们是否相同。我们将使用 Python 的time函数来计算从开始到结束的所有操作的持续时间：

data = []
data_gpu = []
gpu_out = []

# generate random arrays.
for _ in range(num_arrays):
    data.append(np.random.randn(array_len).astype('float32'))

t_start = time()

# copy arrays to GPU.
for k in range(num_arrays):
    data_gpu.append(gpuarray.to_gpu(data[k]))

# process arrays.
for k in range(num_arrays):
    mult_ker(data_gpu[k], np.int32(array_len), block=(64,1,1), grid=(1,1,1))

# copy arrays from GPU.
for k in range(num_arrays):
    gpu_out.append(data_gpu[k].get())

t_end = time()

for k in range(num_arrays):
    assert (np.allclose(gpu_out[k], data[k]))

print 'Total time: %f' % (t_end - t_start)

我们现在准备运行这个程序。我现在就要运行它：

所以，这个程序完成需要了将近三秒的时间。我们将进行一些简单的修改，使我们的程序可以使用流，然后看看我们是否可以获得任何性能提升（这可以在存储库中的multi-kernel_streams.py文件中看到）。

首先，我们注意到对于每个内核启动，我们有一个单独的数据数组进行处理，这些数组存储在 Python 列表中。我们将不得不为每个单独的数组/内核启动对创建一个单独的流对象，所以让我们首先添加一个空列表，名为streams，用来保存我们的流对象：

data = []
data_gpu = []
gpu_out = []
streams = []

我们现在可以生成一系列流，我们将使用它们来组织内核启动。我们可以使用pycuda.driver子模块和Stream类来获取流对象。由于我们已经导入了这个子模块并将其别名为drv，我们可以用以下方式用新的流对象填充我们的列表：

for _ in range(num_arrays):
    streams.append(drv.Stream())

现在，我们将首先修改将数据传输到 GPU 的内存操作。考虑以下步骤：

1. 查找第一个循环，使用gpuarray.to_gpu函数将数组复制到 GPU。我们将希望切换到此函数的异步和适用于流的版本gpu_array.to_gpu_async。（现在我们还必须使用stream参数指定每个内存操作应该使用哪个流）：

for k in range(num_arrays):
    data_gpu.append(gpuarray.to_gpu_async(data[k], stream=streams[k]))

1. 我们现在可以启动我们的内核。这与以前完全相同，只是我们必须通过使用stream参数来指定要使用哪个流：

for k in range(num_arrays):
    mult_ker(data_gpu[k], np.int32(array_len), block=(64,1,1), grid=(1,1,1), stream=streams[k])

1. 最后，我们需要从 GPU 中取出数据。我们可以通过将gpuarray get函数切换为get_async来实现这一点，并再次使用stream参数，如下所示：

for k in range(num_arrays):
    gpu_out.append(data_gpu[k].get_async(stream=streams[k]))

我们现在准备运行我们适用于流的修改后的程序：

在这种情况下，我们获得了三倍的性能提升，考虑到我们需要进行的修改非常少，这并不算太糟糕。但在我们继续之前，让我们尝试更深入地理解为什么这样做有效。

让我们考虑两个 CUDA 内核启动的情况。在我们启动内核之前和之后，我们还将执行与每个内核相对应的 GPU 内存操作，总共有六个操作。我们可以通过图表来可视化在 GPU 上随时间发生的操作，如此—在x轴上向右移动对应于时间持续时间，而y轴对应于在特定时间执行的 GPU 上的操作。这可以用以下图表来描述：

很容易想象为什么流在性能提高方面效果如此好——因为单个流中的操作直到所有必要的先前操作完成之后才会被阻塞，我们将获得不同 GPU 操作之间的并发性，并充分利用我们的设备。这可以通过并发操作的大量重叠来看出。我们可以将基于流的并发性随时间的变化可视化如下：

使用 CUDA 流进行并发康威生命游戏

我们现在将看到一个更有趣的应用——我们将修改上一章的 LIFE（康威的生命游戏）模拟，以便我们将同时显示四个独立的动画窗口。（建议您查看上一章的这个示例，如果您还没有的话。）

让我们从上一章的存储库中获取旧的 LIFE 模拟，应该在4目录中的conway_gpu.py下。现在我们将把这个修改为基于新的 CUDA 流的并发 LIFE 模拟。（我们将在稍后看到的这种基于流的新模拟也可以在本章目录5中的conway_gpu_streams.py文件中找到。）

转到文件末尾的主函数。我们将设置一个新变量，用num_concurrent表示我们将同时显示多少个并发动画（其中N表示模拟栅格的高度/宽度，与以前一样）。我们在这里将其设置为4，但您可以随意尝试其他值：

if __name__ == '__main__':

    N = 128
    num_concurrent = 4

我们现在需要一组num_concurrent流对象，并且还需要在 GPU 上分配一组输入和输出栅格。当然，我们只需将这些存储在列表中，并像以前一样初始化栅格。我们将设置一些空列表，并通过循环填充每个适当的对象，如下所示（请注意我们如何在每次迭代中设置一个新的初始状态栅格，将其发送到 GPU，并将其连接到lattices_gpu）：

streams = []
lattices_gpu = []
newLattices_gpu = []

for k in range(num_concurrent):
    streams.append(drv.Stream())
    lattice = np.int32( np.random.choice([1,0], N*N, p=[0.25, 0.75]).reshape(N, N) )
    lattices_gpu.append(gpuarray.to_gpu(lattice)) 
    newLattices_gpu.append(gpuarray.empty_like(lattices_gpu[k])) 

由于我们只在程序启动期间执行此循环一次，并且几乎所有的计算工作都将在动画循环中进行，因此我们实际上不必担心实际使用我们刚刚生成的流。

我们现在将使用 Matplotlib 使用子图函数设置环境；请注意，我们可以通过设置ncols参数来设置多个动画图。我们将有另一个列表结构，它将对应于动画更新中所需的图像imgs。请注意，我们现在可以使用get_async和相应的流来设置这个：

fig, ax = plt.subplots(nrows=1, ncols=num_concurrent)
imgs = []

for k in range(num_concurrent):
    imgs.append( ax[k].imshow(lattices_gpu[k].get_async(stream=streams[k]), interpolation='nearest') )

在主函数中要更改的最后一件事是以ani = animation.FuncAnimation开头的倒数第二行。让我们修改update_gpu函数的参数，以反映我们正在使用的新列表，并添加两个参数，一个用于传递我们的streams列表，另一个用于指示应该有多少并发动画的参数：

ani = animation.FuncAnimation(fig, update_gpu, fargs=(imgs, newLattices_gpu, lattices_gpu, N, streams, num_concurrent) , interval=0, frames=1000, save_count=1000)    

现在我们需要对update_gpu函数进行必要的修改以接受这些额外的参数。在文件中向上滚动一点，并按照以下方式修改参数：

def update_gpu(frameNum, imgs, newLattices_gpu, lattices_gpu, N, streams, num_concurrent):

现在，我们需要修改这个函数，使其迭代num_concurrent次，并像以前一样设置imgs的每个元素，最后返回整个imgs列表：

for k in range(num_concurrent):
    conway_ker( newLattices_gpu[k], lattices_gpu[k], grid=(N/32,N/32,1), block=(32,32,1), stream=streams[k] )
     imgs[k].set_data(newLattices_gpu[k].get_async(stream=streams[k]) )
     lattices_gpu[k].set_async(newLattices_gpu[k], stream=streams[k])

 return imgs

注意我们所做的更改——每个内核都在适当的流中启动，而get已经切换为与相同流同步的get_async。

最后，在循环中的最后一行将 GPU 数据从一个设备数组复制到另一个设备数组，而不进行任何重新分配。在此之前，我们可以使用简写切片操作符[:]直接在数组之间复制元素，而不在 GPU 上重新分配任何内存；在这种情况下，切片操作符符号充当 PyCUDA set函数的别名，用于 GPU 数组。 （当然，set是将一个 GPU 数组复制到另一个相同大小的数组，而不进行任何重新分配的函数。）幸运的是，确实有一个流同步版本的这个函数，set_async，但我们需要明确地使用这个函数来调用这个函数，明确指定要复制的数组和要使用的流。

我们现在已经完成并准备运行。转到终端并在命令行输入python conway_gpu_streams.py来观看展示：

事件

事件是存在于 GPU 上的对象，其目的是作为操作流的里程碑或进度标记。事件通常用于在设备端精确计时操作的时间持续时间；到目前为止，我们所做的测量都是使用基于主机的 Python 分析器和标准 Python 库函数，如time。此外，事件还可以用于向主机提供有关流状态和已完成的操作的状态更新，以及用于显式基于流的同步。

让我们从一个不使用显式流并使用事件来测量仅一个单个内核启动的示例开始。（如果我们的代码中没有显式使用流，CUDA 实际上会隐式定义一个默认流，所有操作都将放入其中）。

在这里，我们将使用与本章开头相同的无用的乘法/除法循环内核和标题，并修改大部分以下内容。我们希望为此示例运行一个长时间的单个内核实例，因此我们将生成一个巨大的随机数数组供内核处理，如下所示：

array_len = 100*1024**2
data = np.random.randn(array_len).astype('float32')
data_gpu = gpuarray.to_gpu(data)

现在，我们使用pycuda.driver.Event构造函数构造我们的事件（当然，pycuda.driver已经被我们之前的导入语句别名为drv）。

我们将在这里创建两个事件对象，一个用于内核启动的开始，另一个用于内核启动的结束（我们将始终需要两个事件对象来测量任何单个 GPU 操作，很快我们将看到）：

start_event = drv.Event()
end_event = drv.Event()

现在，我们准备启动我们的内核，但首先，我们必须使用事件记录函数标记start_event实例在执行流中的位置。我们启动内核，然后标记end_event在执行流中的位置，并且使用record：

start_event.record()
mult_ker(data_gpu, np.int32(array_len), block=(64,1,1), grid=(1,1,1))
end_event.record()

事件具有一个二进制值，指示它们是否已经到达，这是由函数 query 给出的。让我们在内核启动后立即为两个事件打印状态更新：

print 'Has the kernel started yet? {}'.format(start_event.query())
 print 'Has the kernel ended yet? {}'.format(end_event.query())

现在让我们运行这个，看看会发生什么：

我们的目标是最终测量内核执行的时间持续，但内核似乎还没有启动。在 PyCUDA 中，内核是异步启动的（无论它们是否存在于特定流中），因此我们必须确保我们的主机代码与 GPU 正确同步。

由于end_event是最后一个，我们可以通过此事件对象的synchronize函数阻止进一步的主机代码执行，直到内核完成；这将确保在执行任何进一步的主机代码之前内核已经完成。让我们在适当的位置添加一行代码来做到这一点：

end_event.synchronize()

print 'Has the kernel started yet?  {}'.format(start_event.query())

print 'Has the kernel ended yet? {}'.format(end_event.query())

最后，我们准备测量内核的执行时间；我们可以使用事件对象的time_till或time_since操作来与另一个事件对象进行比较，以获取这两个事件之间的时间（以毫秒为单位）。让我们使用start_event的time_till操作在end_event上：

print 'Kernel execution time in milliseconds: %f ' % start_event.time_till(end_event)

时间持续可以通过 GPU 上已经发生的两个事件之间的time_till和time_since函数来测量。请注意，这些函数始终以毫秒为单位返回值！

现在让我们再次运行我们的程序：

（此示例也可在存储库中的simple_event_example.py文件中找到。）

事件和流

我们现在将看到如何在流方面使用事件对象；这将使我们对各种 GPU 操作的流程具有高度复杂的控制，使我们能够准确了解每个单独流的进展情况，甚至允许我们在忽略其他流的情况下与主机同步特定流。

首先，我们必须意识到这一点——每个流必须有自己专用的事件对象集合；多个流不能共享一个事件对象。让我们通过修改之前的示例multi_kernel_streams.py来确切了解这意味着什么。在内核定义之后，让我们添加两个额外的空列表——start_events和end_events。我们将用事件对象填充这些列表，这些事件对象将对应于我们拥有的每个流。这将允许我们在每个流中计时一个 GPU 操作，因为每个 GPU 操作都需要两个事件：

data = []
data_gpu = []
gpu_out = []
streams = []
start_events = []
end_events = []

for _ in range(num_arrays):
    streams.append(drv.Stream())
    start_events.append(drv.Event())
    end_events.append(drv.Event())

现在，我们可以通过修改第二个循环来逐个计时每个内核启动，使用事件的开始和结束记录。请注意，由于这里有多个流，我们必须将适当的流作为参数输入到每个事件对象的record函数中。还要注意，我们可以在第二个循环中捕获结束事件；这仍然可以完美地捕获内核执行持续时间，而不会延迟启动后续内核。现在考虑以下代码：

for k in range(num_arrays):
    start_events[k].record(streams[k])
    mult_ker(data_gpu[k], np.int32(array_len), block=(64,1,1), grid=(1,1,1), stream=streams[k])

for k in range(num_arrays):
    end_events[k].record(streams[k])

现在我们将提取每个单独内核启动的持续时间。让我们在迭代断言检查之后添加一个新的空列表，并通过time_till函数将其填充为持续时间：

kernel_times = []
for k in range(num_arrays):
   kernel_times.append(start_events[k].time_till(end_events[k]))

现在让我们在最后添加两个print语句，告诉我们内核执行时间的平均值和标准偏差：

print 'Mean kernel duration (milliseconds): %f' % np.mean(kernel_times)
print 'Mean kernel standard deviation (milliseconds): %f' % np.std(kernel_times)

我们现在可以运行这个：

（此示例也可在存储库中的multi-kernel_events.py中找到。）

我们看到内核持续时间的标准偏差相对较低，这是好事，考虑到每个内核在相同的块和网格大小上处理相同数量的数据——如果存在较高的偏差，那将意味着我们在内核执行中使用 GPU 的不均匀程度很高，我们将不得不重新调整参数以获得更高的并发水平。

上下文

CUDA 上下文通常被描述为类似于操作系统中的进程。让我们来看看这意味着什么——进程是计算机上运行的单个程序的实例；操作系统内核之外的所有程序都在一个进程中运行。每个进程都有自己的一组指令、变量和分配的内存，一般来说，它对其他进程的操作和内存是盲目的。当一个进程结束时，操作系统内核会进行清理，确保进程分配的所有内存都已被释放，并关闭进程使用的任何文件、网络连接或其他资源。（好奇的 Linux 用户可以使用命令行top命令查看计算机上运行的进程，而 Windows 用户可以使用 Windows 任务管理器查看）。

类似于进程，上下文与正在使用 GPU 的单个主机程序相关联。上下文在内存中保存了所有正在使用的 CUDA 内核和分配的内存，并且对其他当前存在的上下文的内核和内存是盲目的。当上下文被销毁（例如，在基于 GPU 的程序结束时），GPU 会清理上下文中的所有代码和分配的内存，为其他当前和未来的上下文释放资源。到目前为止，我们编写的程序都存在于一个单一的上下文中，因此这些操作和概念对我们来说是不可见的。

让我们也记住，一个单独的程序开始时是一个单独的进程，但它可以分叉自身以在多个进程或线程中运行。类似地，一个单独的 CUDA 主机程序可以在 GPU 上生成和使用多个 CUDA 上下文。通常，当我们分叉主机进程的新进程或线程时，我们将创建一个新的上下文以获得主机端的并发性。（然而，应该强调的是，主机进程和 CUDA 上下文之间没有确切的一对一关系）。

与生活中的许多其他领域一样，我们将从一个简单的例子开始。我们将首先看看如何访问程序的默认上下文并在其间进行同步。

同步当前上下文

我们将看到如何在 Python 中显式同步我们的设备上下文，就像在 CUDA C 中一样；这实际上是 CUDA C 中最基本的技能之一，在其他大多数关于这个主题的书籍的第一章或第二章中都有涵盖。到目前为止，我们已经能够避免这个话题，因为 PyCUDA 自动使用pycuda.gpuarray函数（如to_gpu或get）执行大多数同步；否则，在本章开头我们看到的to_gpu_async或get_async函数的情况下，同步是由流处理的。

我们将谦卑地开始修改我们在第三章中编写的程序，使用 PyCUDA 入门，它使用显式上下文同步生成 Mandelbrot 集的图像。（这在存储库的3目录下的文件gpu_mandelbrot0.py中可用。）

我们在这里不会获得任何性能提升，与我们最初的 Mandelbrot 程序相比；这个练习的唯一目的是帮助我们理解 CUDA 上下文和 GPU 同步。

看一下头部，我们当然看到了import pycuda.autoinit一行。我们可以使用pycuda.autoinit.context访问当前上下文对象，并且可以通过调用pycuda.autoinit.context.synchronize()函数在当前上下文中进行同步。

现在让我们修改gpu_mandelbrot函数以处理显式同步。我们看到的第一行与 GPU 相关的代码是这样的：

mandelbrot_lattice_gpu = gpuarray.to_gpu(mandelbrot_lattice)

我们现在可以将其更改为显式同步。我们可以使用to_gpu_async异步将数据复制到 GPU，然后进行同步，如下所示：

mandelbrot_lattice_gpu = gpuarray.to_gpu_async(mandelbrot_lattice) pycuda.autoinit.context.synchronize()

然后我们看到下一行使用gpuarray.empty函数在 GPU 上分配内存。由于 GPU 架构的性质，CUDA 中的内存分配始终是自动同步的；这里没有异步内存分配的等价物。因此，我们保持这行与之前一样。

CUDA 中的内存分配始终是同步的！

现在我们看到接下来的两行 - 我们的 Mandelbrot 内核通过调用mandel_ker启动，并且我们通过调用get复制了 Mandelbrot gpuarray对象的内容。在内核启动后我们进行同步，将get切换为get_async，最后进行最后一行同步：

mandel_ker( mandelbrot_lattice_gpu, mandelbrot_graph_gpu, np.int32(max_iters), np.float32(upper_bound))
pycuda.autoinit.context.synchronize()
mandelbrot_graph = mandelbrot_graph_gpu.get_async()
pycuda.autoinit.context.synchronize()

现在我们可以运行这个程序，它将像第三章中的Getting Started with PyCUDA一样在磁盘上生成一个 Mandelbrot 图像。

（此示例也可在存储库中的gpu_mandelbrot_context_sync.py中找到。）

手动上下文创建

到目前为止，我们一直在所有 PyCUDA 程序的开头导入pycuda.autoinit；这实际上在程序开始时创建一个上下文，并在结束时销毁它。

让我们尝试手动操作。我们将制作一个小程序，只是将一个小数组复制到 GPU，然后将其复制回主机，打印数组，然后退出。

我们从导入开始：

import numpy as np
from pycuda import gpuarray
import pycuda.driver as drv

首先，我们使用pycuda.driver.init函数初始化 CUDA，这里被别名为drv：

drv.init()

现在我们选择要使用的 GPU；在拥有多个 GPU 的情况下是必要的。我们可以使用pycuda.driver.Device选择特定的 GPU；如果您只有一个 GPU，就像我一样，可以使用pycuda.driver.Device(0)访问它，如下所示：

dev = drv.Device(0)

现在我们可以使用make_context在此设备上创建一个新上下文，如下所示：

ctx = dev.make_context()

现在我们有了一个新的上下文，这将自动成为默认上下文。让我们将一个数组复制到 GPU，然后将其复制回主机并打印出来：

x = gpuarray.to_gpu(np.float32([1,2,3]))
print x.get()

现在我们完成了。我们可以通过调用pop函数销毁上下文：

ctx.pop()

就是这样！我们应该始终记得在程序退出之前销毁我们显式创建的上下文。

（此示例可以在存储库中的本章目录下的simple_context_create.py文件中找到。）

主机端多进程和多线程

当然，我们可以通过在主机的 CPU 上使用多个进程或线程来实现并发。让我们现在明确区分主机端操作系统进程和线程。

每个存在于操作系统内核之外的主机端程序都作为一个进程执行，并且也可以存在于多个进程中。进程有自己的地址空间，因为它与所有其他进程同时运行并独立运行。进程通常对其他进程的操作视而不见，尽管多个进程可以通过套接字或管道进行通信。在 Linux 和 Unix 中，使用 fork 系统调用生成新进程。

相比之下，主机端线程存在于单个进程中，多个线程也可以存在于单个进程中。单个进程中的多个线程并发运行。同一进程中的所有线程共享进程内的相同地址空间，并且可以访问相同的共享变量和数据。通常，资源锁用于在多个线程之间访问数据，以避免竞争条件。在编译语言（如 C、C++或 Fortran）中，多个进程线程通常使用 Pthreads 或 OpenMP API 进行管理。

线程比进程更轻量级，操作系统内核在单个进程中的多个线程之间切换任务要比在多个进程之间切换任务快得多。通常，操作系统内核会自动在不同的 CPU 核心上执行不同的线程和进程，以建立真正的并发性。

Python 的一个特点是，虽然它通过threading模块支持多线程，但所有线程都将在同一个 CPU 核心上执行。这是由于 Python 是一种解释脚本语言的技术细节所致，与 Python 的全局标识符锁（GIL）有关。要通过 Python 在主机上实现真正的多核并发，不幸的是，我们必须使用multiprocessing模块生成多个进程。（不幸的是，由于 Windows 处理进程的方式，multiprocessing模块目前在 Windows 下并不完全可用。Windows 用户将不幸地不得不坚持单核多线程，如果他们想要在主机端实现任何形式的并发。）

我们现在将看到如何在 Python 中使用两个线程来使用基于 GPU 的操作；Linux 用户应该注意，通过将threading的引用切换到multiprocessing，将Thread的引用切换到Process，可以很容易地将其扩展到进程，因为这两个模块看起来和行为都很相似。然而，由于 PyCUDA 的性质，我们将不得不为每个将使用 GPU 的线程或进程创建一个新的 CUDA 上下文。让我们立即看看如何做到这一点。

主机并发的多个上下文

让我们首先简要回顾一下如何在 Python 中创建一个可以通过简单示例返回值给主机的单个主机线程。（这个例子也可以在存储库中的single_thread_example.py文件的5下看到。）我们将使用threading模块中的Thread类来创建Thread的子类，如下所示：

import threading
class PointlessExampleThread(threading.Thread):

现在我们设置我们的构造函数。我们调用父类的构造函数，并在对象中设置一个空变量，这将是线程返回的值：

def __init__(self):
    threading.Thread.__init__(self)
    self.return_value = None

现在我们在我们的线程类中设置运行函数，这是在启动线程时将被执行的内容。我们只需要让它打印一行并设置返回值：

def run(self):
    print 'Hello from the thread you just spawned!'
    self.return_value = 123

最后，我们需要设置 join 函数。这将允许我们从线程中接收一个返回值：

def join(self):
    threading.Thread.join(self)
    return self.return_value

现在我们已经设置好了我们的线程类。让我们将这个类的一个实例作为NewThread对象启动，通过调用start方法来生成新的线程，然后通过调用join来阻塞执行并从主机线程获取输出：

NewThread = PointlessExampleThread()
NewThread.start()
thread_output = NewThread.join()
print 'The thread completed and returned this value: %s' % thread_output

现在让我们运行这个：

现在，我们可以在主机上的多个并发线程之间扩展这个想法，通过多个上下文和线程来启动并发的 CUDA 操作。现在我们来看最后一个例子。让我们重新使用本章开头的无意义的乘法/除法内核，并在我们生成的每个线程中启动它。

首先，让我们看一下导入部分。由于我们正在创建显式上下文，请记住删除pycuda.autoinit，并在最后添加一个threading导入：

import pycuda
import pycuda.driver as drv
from pycuda import gpuarray
from pycuda.compiler import SourceModule
import numpy as np
from time import time
import threading 

我们将使用与之前相同的数组大小，但这次线程的数量和数组的数量将直接对应。通常情况下，我们不希望在主机上生成超过 20 个左右的线程，所以我们只会使用10个数组。因此，现在考虑以下代码：

num_arrays = 10
array_len = 1024**2

现在，我们将把旧的内核存储为一个字符串对象；由于这只能在一个上下文中编译，我们将不得不在每个线程中单独编译这个内核：

kernel_code = """ 
__global__ void mult_ker(float * array, int array_len)
{
     int thd = blockIdx.x*blockDim.x + threadIdx.x;
     int num_iters = array_len / blockDim.x;
    for(int j=0; j < num_iters; j++)
     {
     int i = j * blockDim.x + thd;
     for(int k = 0; k < 50; k++)
     {
         array[i] *= 2.0;
         array[i] /= 2.0;
     }
 }
}
"""

现在我们可以开始设置我们的类。我们将像以前一样创建threading.Thread的另一个子类，并设置构造函数以接受一个输入数组作为参数。我们将用None初始化一个输出变量，就像以前一样：

class KernelLauncherThread(threading.Thread):
    def __init__(self, input_array):
        threading.Thread.__init__(self)
        self.input_array = input_array
        self.output_array = None

现在我们可以编写run函数。我们选择我们的设备，在该设备上创建一个上下文，编译我们的内核，并提取内核函数引用。注意self对象的使用：

def run(self):
    self.dev = drv.Device(0)
    self.context = self.dev.make_context()
    self.ker = SourceModule(kernel_code)
    self.mult_ker = self.ker.get_function('mult_ker')

现在我们将数组复制到 GPU，启动内核，然后将输出复制回主机。然后我们销毁上下文：

self.array_gpu = gpuarray.to_gpu(self.input_array)
self.mult_ker(self.array_gpu, np.int32(array_len), block=(64,1,1), grid=(1,1,1))
self.output_array = self.array_gpu.get()
self.context.pop()

最后，我们设置了 join 函数。这将把output_array返回到主机：

 def join(self):
     threading.Thread.join(self)
     return self.output_array

我们现在已经完成了我们的子类。我们将设置一些空列表来保存我们的随机测试数据、线程对象和线程输出值，与之前类似。然后我们将生成一些随机数组进行处理，并设置一个核发射器线程列表，这些线程将分别操作每个数组：

data = []
gpu_out = []
threads = []
for _ in range(num_arrays):
    data.append(np.random.randn(array_len).astype('float32'))
for k in range(num_arrays):
 threads.append(KernelLauncherThread(data[k]))

现在我们将启动每个线程对象，并通过使用join将其输出提取到gpu_out列表中：

for k in range(num_arrays):
    threads[k].start()

for k in range(num_arrays):
    gpu_out.append(threads[k].join())

最后，我们只需对输出数组进行简单的断言，以确保它们与输入相同：

for k in range(num_arrays):
    assert (np.allclose(gpu_out[k], data[k]))

这个示例可以在存储库中的multi-kernel_multi-thread.py文件中看到。

总结

我们通过学习设备同步和从主机上同步 GPU 操作的重要性开始了本章；这允许依赖操作在进行之前完成。到目前为止，这个概念一直被 PyCUDA 自动处理，然后我们学习了 CUDA 流，它允许在 GPU 上同时执行独立的操作序列而无需在整个 GPU 上进行同步，这可以大大提高性能；然后我们学习了 CUDA 事件，它允许我们在给定流中计时单个 CUDA 内核，并确定流中的特定操作是否已发生。接下来，我们学习了上下文，它类似于主机操作系统中的进程。我们学习了如何在整个 CUDA 上下文中显式同步，然后看到了如何创建和销毁上下文。最后，我们看到了如何在 GPU 上生成多个上下文，以允许主机上的多个线程或进程使用 GPU。

问题

1. 在第一个示例中内核的启动参数中，我们的内核每个都是在 64 个线程上启动的。如果我们将线程数增加到并超过 GPU 中的核心数，这会如何影响原始版本和流版本的性能？

2. 考虑本章开头给出的 CUDA C 示例，它演示了cudaDeviceSynchronize的使用。您认为在不使用流，只使用cudaDeviceSynchronize的情况下，是否可能在多个内核之间获得一定程度的并发？

3. 如果您是 Linux 用户，请修改上一个示例，使其在进程上运行而不是线程。

4. 考虑multi-kernel_events.py程序；我们说内核执行持续时间的标准偏差很低是件好事。如果标准偏差很高会有什么坏处？

5. 在上一个示例中，我们只使用了 10 个主机端线程。列出两个原因，说明为什么我们必须使用相对较少数量的线程或进程来在主机上启动并发的 GPU 操作。

第六章：调试和分析您的 CUDA 代码

在本章中，我们将最终学习如何使用多种不同的方法和工具调试和分析我们的 GPU 代码。虽然我们可以使用 Spyder 和 PyCharm 等 IDE 轻松调试纯 Python 代码，但我们无法使用这些工具来调试实际的 GPU 代码，记住 GPU 代码本身是用 CUDA-C 编写的，PyCUDA 提供了一个接口。调试 CUDA 内核的第一种最简单的方法是使用printf语句，我们实际上可以直接在 CUDA 内核中调用它来打印到标准输出。我们将看到如何在 CUDA 的上下文中使用printf以及如何有效地应用它进行调试。

接下来，我们将填补 CUDA-C 编程中的一些空白，以便我们可以直接在 NVIDIA Nsight IDE 中编写 CUDA 程序，这将允许我们为我们一直在编写的一些代码创建 CUDA-C 的测试用例。我们将看看如何使用nvcc命令行编译 CUDA-C 程序，以及如何在 Nsight IDE 中进行编译。然后，我们将看看如何在 Nsight 中进行调试，并使用 Nsight 了解 CUDA lockstep 属性。最后，我们将概述 NVIDIA 命令行和 Visual Profilers 以对我们的代码进行分析。

本章的学习成果包括以下内容：

• 有效地使用printf作为 CUDA 内核的调试工具

• 在 Python 之外编写完整的 CUDA-C 程序，特别是用于创建调试的测试用例

• 使用nvcc编译器在命令行上编译 CUDA-C 程序

• 使用 NVIDIA Nsight IDE 开发和调试 CUDA 程序

• 了解 CUDA warp lockstep 属性以及为什么我们应该避免单个 CUDA warp 内的分支分歧

• 学会有效使用 NVIDIA 命令行和 Visual Profilers 进行 GPU 代码的调试

技术要求

本章需要一台安装了现代 NVIDIA GPU（2016 年以后）的 Linux 或 Windows 10 PC，并安装了所有必要的 GPU 驱动程序和 CUDA Toolkit（9.0 及以上）。还需要一个合适的 Python 2.7 安装（如 Anaconda Python 2.7），并安装了 PyCUDA 模块。

本章的代码也可以在 GitHub 上找到：github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA。

有关先决条件的更多信息，请查看本书的前言，有关软件和硬件要求，请查看github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA中的 README。

在 CUDA 内核中使用 printf

也许会让人惊讶，但我们实际上可以直接从 CUDA 内核中将文本打印到标准输出；不仅如此，每个单独的线程都可以打印自己的输出。当我们调试内核时，这将特别方便，因为我们可能需要监视特定变量或代码中特定点的计算值，这也将使我们摆脱使用调试器逐步进行调试的束缚。从 CUDA 内核中打印输出的方法是使用 C/C++编程中最基本的函数，大多数人在编写他们的第一个 C 程序Hello world时会学到的函数：printf。当然，printf是将字符串打印到标准输出的标准函数，实际上在 C 编程语言中相当于 Python 的print函数。

现在让我们简要回顾一下如何在 CUDA 中使用printf。首先要记住的是，printf总是以字符串作为其第一个参数；因此，在 C 中打印"Hello world!"是用printf("Hello world!\n");来完成的。（当然，\n表示"新行"或"返回"，它将输出在终端上移到下一行。）printf还可以在我们想要直接在 C 中打印任何常量或变量的情况下，采用可变数量的参数：如果我们想要将123整数打印到输出，我们可以使用printf("%d", 123);（其中%d表示字符串后面跟着一个整数。）

类似地，我们使用%f，%e或%g来打印浮点值（其中%f是十进制表示法，%e是科学表示法，%g是最短的表示法，无论是十进制还是科学表示法）。我们甚至可以连续打印几个值，记得按正确的顺序放置这些指示符：printf("%d is a prime number, %f is close to pi, and %d is even.\n", 17, 3.14, 4);将在终端上打印"17 is a prime number, 3.14 is close to pi, and 4 is even."。

现在，在这本书的近一半时，我们终于要开始创建我们的第一个并行Hello world程序了！我们首先导入适当的模块到 Python 中，然后编写我们的内核。我们将首先打印每个单独线程的线程和网格标识（我们只会在一维块和网格中启动这个，所以我们只需要x值）：

ker = SourceModule('''
__global__ void hello_world_ker()
{
    printf("Hello world from thread %d, in block %d!\\n", threadIdx.x, blockIdx.x);

让我们停下来注意一下，我们写的是\\n而不是\n。这是因为 Python 中的三引号本身会将\n解释为"新行"，所以我们必须使用双反斜杠来表示我们是字面意思，以便将\n直接传递给 CUDA 编译器。

我们现在将打印有关块和网格维度的一些信息，但我们希望确保它在每个线程完成其初始printf命令后打印。我们可以通过放入__syncthreads();来确保每个单独的线程在执行第一个printf函数后同步。

现在，我们只想将块和网格的维度打印到终端一次；如果我们在这里放置printf语句，每个线程都会打印相同的信息。我们可以通过只有一个指定的线程打印输出来实现这一点；让我们选择第 0 个块的第 0 个线程，这是唯一保证存在的线程，无论我们选择的块和网格的维度如何。我们可以通过 C 的if语句来实现这一点：

 if(threadIdx.x == 0 && blockIdx.x == 0)
 {

我们现在将打印出我们的块和网格的维度，并关闭if语句，这将是我们的 CUDA 内核的结束：

 printf("-------------------------------------\\n");
 printf("This kernel was launched over a grid consisting of %d blocks,\\n", gridDim.x);
 printf("where each block has %d threads.\\n", blockDim.x);
 }
}
''')

我们现在将提取内核，然后在由两个块组成的网格上启动它，每个块有五个线程：

hello_ker = ker.get_function("hello_world_ker")
hello_ker( block=(5,1,1), grid=(2,1,1) )

让我们现在运行这个程序（该程序也可以在存储库中的hello-world_gpu.py的6下找到）：

使用 printf 进行调试

让我们通过一个例子来看看如何使用printf调试 CUDA 内核，然后再继续。这种方法并没有确切的科学依据，但通过经验可以学会这种技能。我们将从一个 CUDA 内核开始，用于矩阵乘法，但其中有几个错误。（鼓励读者随着我们的步骤阅读代码，该代码可在存储库中的6目录中的broken_matrix_ker.py文件中找到。）

在继续之前，让我们简要回顾一下矩阵乘法。假设我们有两个矩阵A和B，我们将它们相乘得到另一个相同大小的矩阵C，如下所示：。我们通过迭代所有元组来做到这一点，并将的值设置为A的第i行和B的第j列的点积：。

换句话说，我们将输出矩阵C中的每个i, j元素设置如下：

假设我们已经编写了一个内核，用于执行矩阵乘法，它接受表示输入矩阵的两个数组，一个额外的预先分配的浮点数组，输出将写入其中，并且一个表示每个矩阵的高度和宽度的整数（我们将假设所有矩阵都是相同大小和正方形的）。这些矩阵都将以一维的float *数组以行优先的一维布局表示。此外，这将被实现为每个 CUDA 线程处理输出矩阵中的单个行/列元组。

我们进行了一个小的测试案例，并将其与 CUDA 中矩阵乘法的输出进行了检查，对于两个 4 x 4 矩阵，它作为断言检查失败，如下所示：

test_a = np.float32( [xrange(1,5)] * 4 )
test_b = np.float32([xrange(14,10, -1)]*4 )
output_mat = np.matmul(test_a, test_b)

test_a_gpu = gpuarray.to_gpu(test_a)
test_b_gpu = gpuarray.to_gpu(test_b)
output_mat_gpu = gpuarray.empty_like(test_a_gpu)

matrix_ker(test_a_gpu, test_b_gpu, output_mat_gpu, np.int32(4), block=(2,2,1), grid=(2,2,1))

assert( np.allclose(output_mat_gpu.get(), output_mat) )

我们现在将运行这个程序，并且不出所料地得到以下输出：

现在让我们来看一下 CUDA C 代码，其中包括一个内核和一个设备函数：

ker = SourceModule('''
// row-column dot-product for matrix multiplication
__device__ float rowcol_dot(float *matrix_a, float *matrix_b, int row, int col, int N)
{
 float val = 0;

 for (int k=0; k < N; k++)
 {
     val += matrix_a[ row + k*N ] * matrix_b[ col*N + k];
 }
 return(val);
}

// matrix multiplication kernel that is parallelized over row/column tuples.

__global__ void matrix_mult_ker(float * matrix_a, float * matrix_b, float * output_matrix, int N)
{
 int row = blockIdx.x + threadIdx.x;
 int col = blockIdx.y + threadIdx.y;

 output_matrix[col + row*N] = rowcol_dot(matrix_a, matrix_b, col, row, N);
}
''')

我们的目标是在我们的 CUDA 代码中聪明地放置printf调用，以便我们可以监视内核和设备函数中的许多适当的值和变量；我们还应该确保在每个printf调用中打印出线程和块号。

让我们从内核的入口点开始。我们看到两个变量row和col，所以我们应该立即检查这些。让我们在设置它们之后立即放上一行代码（因为这是在两个维度上并行化的，我们应该打印threadIdx和blockIdx的x和y值）：

printf("threadIdx.x,y: %d,%d blockIdx.x,y: %d,%d -- row is %d, col is %d.\\n", threadIdx.x, threadIdx.y, blockIdx.x, blockIdx.y, row, col);

再次运行代码，我们得到了这个输出：

有两件事情立即引人注目：行和列元组有重复的值（每个单独的元组应该只表示一次），而且行和列的值从未超过两，而它们都应该达到三（因为这个单元测试使用 4 x 4 的矩阵）。这应该告诉我们，我们正在错误地计算行和列的值；确实，我们忘记了将blockIdx的值乘以blockDim的值来找到目标行/列值。我们按照以下方式修复这个问题：

int row = blockIdx.x*blockDim.x + threadIdx.x;
int col = blockIdx.y*blockDim.y + threadIdx.y;

然而，如果我们再次运行程序，我们仍然会得到一个断言错误。让我们保留原始的printf调用，这样我们就可以在继续进行的过程中监视这些值。我们看到在内核中有一个对设备函数rowcol_dot的调用，所以我们决定去看一下。让我们首先确保变量被正确传递到设备函数中，通过在开始处放置这个printf调用：

printf("threadIdx.x,y: %d,%d blockIdx.x,y: %d,%d -- row is %d, col is %d, N is %d.\\n", threadIdx.x, threadIdx.y, blockIdx.x, blockIdx.y, row, col, N);

当我们运行程序时，会有更多的行输出，但我们会看到一行说—threadIdx.x,y: 0,0 blockIdx.x,y: 1,0 -- row is 2, col is 0.，还有另一行说—threadIdx.x,y: 0,0 blockIdx.x,y: 1,0 -- row is 0, col is 2, N is 4。通过threadIdx和blockIdx的值，我们看到这是同一个块中的同一个线程，但row和col的值是颠倒的。实际上，当我们查看rowcol_dot设备函数的调用时，我们看到row和col确实是与设备函数声明中的相反。我们修复了这个问题，但当我们再次运行程序时，又出现了另一个断言错误。

让我们在设备函数中的for循环内放置另一个printf调用；这当然是点积，用于在矩阵A的行与矩阵B的列之间执行点积。我们将检查我们正在相乘的矩阵的值，以及k；我们还将只查看第一个线程的值，否则我们将得到一个不连贯的输出混乱。

if(threadIdx.x == 0 && threadIdx.y == 0 && blockIdx.x == 0 && blockIdx.y == 0)
            printf("Dot-product loop: k value is %d, matrix_a value is %f, matrix_b is %f.\\n", k, matrix_a[ row + k*N ], matrix_b[ col*N + k]);

在继续之前，让我们看一下为我们的单元测试设置的A和B矩阵的值：

我们看到，当我们在列之间切换时，两个矩阵都会变化，但在行之间变化时是恒定的。因此，根据矩阵乘法的性质，矩阵A的值应该在我们的for循环中随着k的变化而变化，而B的值应该保持恒定。让我们再次运行程序并检查相关的输出：

因此，看起来我们没有以正确的方式访问矩阵的元素；记住这些矩阵是以行方式存储的，我们修改索引，以便以正确的方式访问它们的值：

val += matrix_a[ row*N + k ] * matrix_b[ col + k*N];

再次运行程序将不会产生断言错误。恭喜，您刚刚使用唯一的printf调试了一个 CUDA 内核！

用 CUDA-C 填补空白

我们现在将介绍如何编写一个完整的 CUDA-C 程序的基础知识。我们将从小处开始，将我们刚刚在上一节中调试的小矩阵乘法测试程序的修复版本翻译成纯 CUDA-C 程序，然后使用 NVIDIA 的nvcc编译器从命令行编译成本机 Windows 或 Linux 可执行文件（我们将在下一节中看到如何使用 NVIDIA 的 Nsight IDE，所以现在我们只使用文本编辑器和命令行）。同样，鼓励读者在我们进行翻译时查看我们正在翻译的 Python 代码，该代码在存储库中作为matrix_ker.py文件可用。

现在，让我们打开我们最喜欢的文本编辑器，创建一个名为matrix_ker.cu的新文件。扩展名将表明这是一个 CUDA-C 程序，可以使用nvcc编译器进行编译。

CUDA-C 程序和库源代码的文件名总是使用.cu文件扩展名。

让我们从头开始——正如 Python 在程序开头使用import关键字导入库一样，我们回忆 C 语言使用#include。在继续之前，我们需要包含一些导入库。

让我们从这些开始：

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

让我们简要地思考一下我们需要这些的原因：cuda_runtime.h是一个头文件，其中包含了我们程序所需的所有特定 CUDA 数据类型、函数和结构的声明。我们需要为我们编写的任何纯 CUDA-C 程序包含这个头文件。stdio.h当然为我们提供了主机的所有标准 I/O 函数，如printf，我们需要stdlib.h来使用主机上的malloc和free动态内存分配函数。

请记住，始终在每个纯 CUDA-C 程序的开头放置#include <cuda_runtime.h>！

现在，在我们继续之前，我们记得我们最终将不得不检查我们的内核输出与正确的已知输出，就像我们在 NumPy 的allclose函数中所做的那样。不幸的是，在 C 中，我们没有像 Python 中的 NumPy 那样的标准或易于使用的数值数学库。往往，如果是一些简单的东西，编写自己的等效函数会更容易，就像在这种情况下一样。这意味着我们现在将明确地制作我们自己的等效于 NumPy 的allclose。我们将这样做：我们将使用 C 中的#define宏来设置一个名为_EPSILON的值，它将作为一个常数来指示输出和期望输出之间的最小值，以便被认为是相同的，我们还将设置一个名为_ABS的宏，它将告诉我们两个数字之间的绝对差异。我们这样做如下：

#define _EPSILON 0.001
#define _ABS(x) ( x > 0.0f ? x : -x )

现在我们可以创建我们自己的allclose版本。这将接受两个浮点指针和一个整数值len，我们循环遍历两个数组并检查它们：如果任何点的差异超过_EPSILON，我们返回-1，否则我们返回 0，表示这两个数组确实匹配。

我们注意到一件事：由于我们使用 CUDA-C，我们在函数定义之前加上__host__，以表明这个函数是打算在 CPU 上运行而不是在 GPU 上运行的：

__host__ int allclose(float *A, float *B, int len)
{

  int returnval = 0;

  for (int i = 0; i < len; i++)
  {
    if ( _ABS(A[i] - B[i]) > _EPSILON )
    {
      returnval = -1;
      break;
    }
  }

  return(returnval);
}

现在我们可以将设备和内核函数剪切并粘贴到这里，就像它们在我们的 Python 版本中出现的那样：

__device__ float rowcol_dot(float *matrix_a, float *matrix_b, int row, int col, int N)
{
  float val = 0;

  for (int k=0; k < N; k++)
  {
        val += matrix_a[ row*N + k ] * matrix_b[ col + k*N];
  }

  return(val);
}

__global__ void matrix_mult_ker(float * matrix_a, float * matrix_b, float * output_matrix, int N)
{

    int row = blockIdx.x*blockDim.x + threadIdx.x;
    int col = blockIdx.y*blockDim.y + threadIdx.y;

  output_matrix[col + row*N] = rowcol_dot(matrix_a, matrix_b, row, col, N);
}

再次，与__host__相比，注意 CUDA 设备函数前面有__device__，而 CUDA 内核前面有__global__。

现在，就像在任何 C 程序中一样，我们需要编写main函数，它将在主机上运行，我们将在其中设置我们的测试案例，并从中显式地启动我们的 CUDA 内核到 GPU 上。再次，与普通的 C 相比，我们将明确指定这也要在 CPU 上运行，并使用__host__：

__host__ int main()
{

我们将要做的第一件事是选择和初始化我们的 GPU。我们可以使用cudaSetDevice来这样做：

cudaSetDevice(0);

cudaSetDevice(0)将选择默认的 GPU。如果您的系统中安装了多个 GPU，您可以选择并使用它们，而不是使用cudaSetDevice(1)，cudaSetDevice(2)等。

现在我们将设置N，就像在 Python 中一样，以指示矩阵的高度/宽度。由于我们的测试案例将只包括 4 x 4 矩阵，我们将其设置为4。由于我们将使用动态分配的数组和指针，我们还必须设置一个值，以指示我们的测试矩阵将需要的字节数。矩阵将由N x N浮点数组成，我们可以使用 C 中的sizeof关键字确定浮点数所需的字节数：

int N = 4;
int num_bytes = sizeof(float)*N*N;

现在我们设置我们的测试矩阵如下；这些将与我们在 Python 测试程序中看到的test_a和test_b矩阵完全对应（请注意我们如何使用h_前缀来表示这些数组存储在主机上，而不是设备上）：

 float h_A[] = { 1.0, 2.0, 3.0, 4.0, \
                 1.0, 2.0, 3.0, 4.0, \
                 1.0, 2.0, 3.0, 4.0, \
                 1.0, 2.0, 3.0, 4.0 };

 float h_B[] = { 14.0, 13.0, 12.0, 11.0, \
                 14.0, 13.0, 12.0, 11.0, \
                 14.0, 13.0, 12.0, 11.0, \
                 14.0, 13.0, 12.0, 11.0 };

现在我们设置另一个数组，它将指示先前测试矩阵的矩阵乘法的预期输出。我们将不得不明确计算这一点，并将这些值放入我们的 C 代码中。最终，我们将在程序结束时将其与 GPU 输出进行比较，但让我们先设置它并将其放在一边：

float h_AxB[] = { 140.0, 130.0, 120.0, 110.0, \
                 140.0, 130.0, 120.0, 110.0, \
                 140.0, 130.0, 120.0, 110.0, \
                 140.0, 130.0, 120.0, 110.0 };

现在我们声明一些指针，这些指针将存在于 GPU 上的数组，并且我们将复制h_A和h_B的值并指向 GPU 的输出。请注意，我们只是使用标准的浮点指针。还要注意前缀d_ - 这是另一个标准的 CUDA-C 约定，表示这些将存在于设备上：

float * d_A;
float * d_B;
float * d_output;

现在，我们将使用cudaMalloc在设备上为d_A和d_B分配一些内存，这几乎与 C 中的malloc相同；这就是 PyCUDA gpuarray函数（如empty或to_gpu）在本书中无形地调用我们分配 GPU 上的内存数组的方式：

cudaMalloc((float **) &d_A, num_bytes);
cudaMalloc((float **) &d_B, num_bytes);

让我们思考一下这是如何工作的：在 C 函数中，我们可以通过在变量前加上一个取地址运算符（&）来获取变量的地址；如果有一个整数x，我们可以用&x来获取它的地址。&x将是一个指向整数的指针，因此它的类型将是int *。我们可以使用这个来将参数的值设置到 C 函数中，而不仅仅使用纯返回值。

由于cudaMalloc通过参数设置指针而不是使用返回值（与常规的malloc相反），我们必须使用取地址运算符，它将是一个指向指针的指针，因为它是一个指向浮点指针的指针，所以我们必须使用括号显式地转换这个值，因为cudaMalloc可以分配任何类型的数组。最后，在第二个参数中，我们必须指示在 GPU 上分配多少字节；我们之前已经设置了num_bytes，以便它是我们需要保存由浮点数组成的 4 x 4 矩阵所需的字节数，所以我们将其插入并继续。

现在我们可以使用两次cudaMemcpy函数将h_A和h_B的值分别复制到d_A和d_B中：

cudaMemcpy(d_A, h_A, num_bytes, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, num_bytes, cudaMemcpyHostToDevice);

cudaMemcpy总是以目标指针作为第一个参数，源指针作为第二个参数，要复制的字节数作为第三个参数，并且还有一个最后的参数。最后一个参数将指示我们是使用cudaMemcpyHostToDevice从主机到 GPU 进行复制，使用cudaMemcpyDeviceToHost从 GPU 到主机进行复制，还是在 GPU 上的两个数组之间进行复制cudaMemcpyDeviceToDevice。

我们现在将分配一个数组来保存我们在 GPU 上进行矩阵乘法的输出，使用cudaMalloc的另一个调用：

cudaMalloc((float **) &d_output, num_bytes);

最后，当我们想要检查内核的输出时，我们将在主机上设置一些存储 GPU 输出的内存。让我们设置一个常规的 C 浮点指针，并使用malloc分配内存，就像我们通常做的那样：

float * h_output;
h_output = (float *) malloc(num_bytes);

现在，我们几乎准备好启动我们的内核。CUDA 使用一个名为dim3的数据结构来指示内核启动的块和网格大小；我们将设置这些，因为我们想要一个 2 x 2 维度的网格和也是 2 x 2 维度的块：

dim3 block(2,2,1);
dim3 grid(2,2,1);

现在我们准备启动我们的内核；我们使用三角形括号来指示 CUDA-C 编译器内核应该启动的块和网格大小：

matrix_mult_ker <<< grid, block >>> (d_A, d_B, d_output, N);

现在，当然，在我们可以将内核的输出复制回主机之前，我们必须确保内核已经执行完毕。我们通过调用cudaDeviceSynchronize来做到这一点，这将阻止主机向 GPU 发出更多命令，直到内核执行完毕：

cudaDeviceSynchronize();

现在我们可以将内核的输出复制到我们在主机上分配的数组中：

cudaMemcpy(h_output, d_output, num_bytes, cudaMemcpyDeviceToHost);

再次，我们同步：

cudaDeviceSynchronize();

在检查输出之前，我们意识到我们不再需要在 GPU 上分配的任何数组。我们通过在每个数组上调用cudaFree来释放这些内存：

cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_output);

我们已经完成了 GPU 的工作，所以我们调用cudaDeviceReset：

cudaDeviceReset();

现在，我们最终检查我们在主机上复制的输出，使用我们在本章开头编写的allclose函数。如果实际输出与预期输出不匹配，我们打印一个错误并返回-1，否则，我们打印匹配并返回 0。然后我们在程序的main函数上放一个闭括号：

if (allclose(h_AxB, h_output, N*N) < 0)
 {
     printf("Error! Output of kernel does not match expected output.\n");
     free(h_output);
     return(-1);
 }
 else
 {
     printf("Success! Output of kernel matches expected output.\n");
     free(h_output);
     return(0);
 }
}

请注意，由于我们已经为h_output分配了内存，在这两种情况下，我们最后调用了标准的 C free 函数。

现在保存我们的文件，并从命令行编译成 Windows 或 Linux 可执行文件，使用nvcc matrix_ker.cu -o matrix_ker。这应该输出一个二进制可执行文件，matrix_ker.exe（在 Windows 中）或matrix_ker（在 Linux 中）。让我们尝试编译和运行它：

恭喜，您刚刚创建了您的第一个纯 CUDA-C 程序！（此示例在存储库中作为matrix_ker.cu在7下可用。）

使用 Nsight IDE 进行 CUDA-C 开发和调试

现在让我们学习如何使用 Nsight IDE 开发 CUDA-C 程序。我们将看到如何导入我们刚刚编写的程序，并在 Nsight 内部进行编译和调试。请注意，由于在 Windows 下它实际上是 Visual Studio IDE 的插件，在 Linux 下是 Eclipse IDE 的插件，因此 Windows 和 Linux 版本的 Nsight 之间存在差异。我们将在接下来的两个子部分中涵盖两者；如果不适用于您的操作系统，请随意跳过。

在 Windows 中使用 Visual Studio 的 Nsight

打开 Visual Studio，点击文件，然后选择新建|项目....会弹出一个窗口，您可以在其中设置项目类型：选择 NVIDIA 下拉项，然后选择 CUDA 9.2：

给项目取一个合适的名称，然后点击确定。在解决方案资源管理器窗口中应该会出现一个项目，其中包含一个简单的预制 CUDA 测试程序，由一个源文件kernel.cu组成，其中包含一个简单的并行加法内核和测试代码。如果您想查看这是否编译和运行，请点击顶部标有本地 Windows 调试器的绿色向右箭头。一个终端应该弹出，显示内核的一些文本输出，然后立即关闭。

如果您在从 Visual Studio 运行后关闭基于 Windows 终端的应用程序时遇到问题，请尝试在主函数的末尾添加getchar();，这将使终端保持打开状态，直到您按下一个键。（或者，您也可以在程序的末尾使用调试器断点。）

现在，让我们添加刚刚编写的 CUDA-C 程序。在解决方案资源管理器窗口中，右键单击kernel.cu，然后单击kernel.cu上的删除。现在，右键单击项目名称，选择添加，然后选择现有项目。现在我们可以选择一个现有文件，找到matrix_ker.cu的路径，并将其添加到项目中。点击 IDE 顶部标有本地 Windows 调试器的绿色箭头，程序应该会在 Windows 终端中再次编译和运行。这就是我们可以在 Visual Studio 中设置和编译完整的 CUDA 程序的全部步骤。

现在让我们看看如何调试我们的 CUDA 内核。让我们首先在代码的入口点matrix_mult_ker处添加一个断点，我们在那里设置了row和col的值。我们可以通过在窗口的行号左侧的灰色列上单击来添加此断点；每个我们添加的断点都应该在那里显示一个红点。（您可以忽略 Visual Studio 编辑器可能在您的代码下方放置的任何红色波浪线；这是因为 CUDA 不是 Visual Studio 的本地语言）：

现在我们可以开始调试。从顶部菜单中选择 Nsight 下拉菜单，然后选择开始 CUDA 调试。这里可能有两个选项，开始 CUDA 调试（Next-Gen）和开始 CUDA 调试（Legacy）。无论选择哪一个都可以，但是根据您的 GPU，可能会在 Next-Gen 上遇到问题；在这种情况下，请选择 Legacy。

您的程序应该启动，并且调试器应该在我们刚刚设置的内核中的断点处停止。让我们按F10跳过这一行，现在看看row变量是否被正确设置。让我们在变量资源管理器中的本地窗口中查看：

通过检查threadIdx和blockIdx的值，我们可以看到我们当前位于网格中的第一个块中的第一个线程；row设置为0，这确实对应于正确的值。现在，让我们检查一些不同线程的row值。为了做到这一点，我们必须在 IDE 中切换线程焦点；我们可以通过单击上面的 Nsight 下拉菜单，然后选择 Windows|CUDA Debug Focus...来实现这一点。应该会出现一个新菜单，允许您选择一个新的线程和块。在菜单中将线程从 0, 0, 0 更改为 1, 0, 0，然后单击确定：

当您再次检查变量时，您应该看到为此线程设置了正确的row值：

简而言之，这就是您在 Visual Studio 中使用 Nsight 进行调试的方法。我们现在已经掌握了如何在 Windows 中使用 Nsight/Visual Studio 调试 CUDA 程序的基础知识，我们可以像调试常规 Windows 程序一样使用所有常规约定（设置断点，启动调试器，继续/恢复，跳过，步入和步出）。主要的区别在于您必须知道如何在 CUDA 线程和块之间切换以检查变量，否则它基本上是一样的。

在 Linux 中使用 Nsight 与 Eclipse

现在我们将看到如何在 Linux 中使用 Nsight。您可以从桌面上选择它打开 Nsight，也可以使用nsight命令从命令行运行它。Nsight IDE 将打开。从 IDE 顶部，单击文件，然后从下拉菜单中选择新建...，然后选择新建 CUDA C/C++项目。将出现一个新窗口，在这里选择 CUDA Runtime 项目。给项目取一个合适的名字，然后点击下一步。您将被提示提供进一步的设置选项，但默认设置对我们的目的来说现在可以工作得很好。（请确保注意这里第三和第四屏幕中源文件和项目路径的位置。）您将进入最终屏幕，在这里您可以按完成来创建项目：

最后，您将在项目视图中看到您的新项目和一些占位代码；从 CUDA 9.2 开始，这将包括一个倒数内核示例。

现在我们可以导入我们的代码。您可以使用 Nsight 中的编辑器删除默认源文件中的所有代码并剪切粘贴，或者您可以手动从项目的源目录中删除文件，手动将matrix_ker.cu文件复制到源目录中，然后选择刷新 Nsight 中的源目录视图，然后按F5。现在可以使用Ctrl + B构建项目，并使用F11运行它。我们程序的输出应该出现在 IDE 的 Console 子窗口中，如下所示：

现在，我们可以在 CUDA 代码中设置断点；让我们在内核的入口点设置一个断点，那里设置了row值。我们将光标放在 Eclipse 编辑器中的该行上，然后按Ctrl + Shift + B进行设置。

现在，我们可以通过按F11（或单击 bug 图标）开始调试。程序应该在main函数的开头暂停，所以按F8继续到第一个断点。您应该在 IDE 中看到我们的 CUDA 内核中的第一行被箭头指向。让我们通过按F6跳过当前行，确保row已经设置。

现在，我们可以轻松地在 CUDA 网格中切换不同的线程和块，以检查它们当前持有的值：从 IDE 顶部，单击窗口下拉菜单，然后单击显示视图，然后选择 CUDA。应该会打开一个显示当前运行内核的窗口，从这里您可以看到此内核正在运行的所有块的列表。

点击第一个，从这里你将能够看到块内运行的所有单个线程：

现在，我们可以通过单击“变量”选项卡来查看与第一个块中的第一个线程对应的变量，这里，row 应该是 0，正如我们所期望的：

现在，我们可以通过再次转到 CUDA 选项卡，选择适当的线程并切换回来，来检查不同线程的值。让我们留在同一个块中，但这次选择线程（1,0,0），再次检查 row 的值：

我们看到 row 的值现在是 1，正如我们所期望的。

现在，我们已经掌握了如何从 Nisight/Eclipse 在 Linux 中调试 CUDA 程序的基础知识，我们可以像调试其他 IDE 中的常规 Linux 程序一样使用所有常规约定（设置断点、启动调试器、继续/恢复、步进、步入和步出）。主要的区别在于我们必须知道如何在 CUDA 线程和块之间切换以检查变量，否则，它基本上是一样的。

使用 Nisight 来理解 CUDA 中的 warp lockstep 属性

我们现在将使用 Nisight 逐步执行一些代码，以帮助我们更好地理解一些 CUDA GPU 架构，以及内核内的分支是如何处理的。这将使我们对如何编写更有效的 CUDA 内核有一些见解。通过分支，我们指的是 GPU 如何处理 CUDA 内核中的if、else或switch等控制流语句。特别是，我们对内核内的分支分歧如何处理感兴趣，即当内核中的一个线程满足成为if语句的条件时，而另一个线程不满足条件并成为else语句时会发生什么：它们是分歧的，因为它们执行不同的代码片段。

让我们写一个小的 CUDA-C 程序作为实验：我们将从一个小内核开始，如果其threadIdx.x值是偶数，则打印一个输出，如果是奇数，则打印另一个输出。然后，我们编写一个main函数，将这个内核启动到由 32 个不同线程组成的一个单一块上：

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

__global__ void divergence_test_ker()
{
    if( threadIdx.x % 2 == 0)
        printf("threadIdx.x %d : This is an even thread.\n", threadIdx.x);
    else
        printf("threadIdx.x %d : This is an odd thread.\n", threadIdx.x);
}

__host__ int main()
{
    cudaSetDevice(0);
    divergence_test_ker<<<1, 32>>>();
    cudaDeviceSynchronize();
    cudaDeviceReset();
}

（此代码也可在存储库中的divergence_test.cu中找到。）

如果我们从命令行编译和运行这个程序，我们可能天真地期望偶数和奇数线程之间会有交错的字符串序列；或者它们可能会随机交错——因为所有线程都是并发运行并且大约在同一时间分支，这是有道理的。

相反，每次我们运行这个程序，我们总是得到这个输出：

所有与偶数线程对应的字符串都先打印出来，而所有与奇数线程对应的字符串都在第二次打印出来。也许 Nisight 调试器可以解释一些问题；让我们像在上一节中那样将这个小程序导入 Nisight 项目，并在内核的第一个if语句处设置断点。然后我们将执行step over，这样调试器就会在第一个printf语句处停下来。由于 Nisight 中的默认线程是（0,0,0），这应该满足了第一个if语句，所以它会一直停在那里，直到调试器继续。

让我们切换到一个奇数线程，比如（1,0,0），看看它现在在我们的程序中的位置：

非常奇怪！线程（1,0,0）在执行中也与线程（0,0,0）处于相同的位置。实际上，如果我们在这里检查每一个其他奇数线程，它们都会停在同一个地方——在一个所有奇数线程应该跳过的printf语句处。

这是什么？这被称为warp 锁步特性。CUDA 架构中的一个warp是一个由 32 个“通道”组成的单元，在这个单元中，我们的 GPU 执行内核和网格，其中每个通道将执行一个线程。warp 的一个主要限制是，所有在一个 warp 上执行的线程必须以锁步的方式执行相同的代码；这意味着并非每个线程确实运行相同的代码，而只是忽略对它不适用的步骤。（这被称为锁步，因为就像一群士兵一起锁步行进一样——不管他们是否想要行进！）

锁步特性意味着如果一个 warp 上运行的单个线程在单个if语句中与其他 31 个线程产生分歧，所有其他 31 个线程的执行都会延迟，直到这个单独的异常线程完成并从其孤立的if分歧中返回。这是在编写内核时应该牢记的一个特性，也是为什么在 CUDA 编程中，分支分歧应该尽量减少的一般规则。

使用 NVIDIA nvprof 分析器和 Visual Profiler

最后，我们将简要概述命令行 Nvidia nvprof分析器。与 Nsight IDE 相比，我们可以自由使用我们编写的任何 Python 代码——我们不必在这里强制编写完整的、纯粹的 CUDA-C 测试函数代码。

我们可以使用nvprof program命令对二进制可执行程序进行基本分析；同样，我们可以使用python命令作为第一个参数，使用脚本作为第二个参数来对 Python 脚本进行分析，如下所示：nvprof python program.py。让我们使用nvprof matrix_ker对我们之前编写的简单矩阵乘法 CUDA-C 可执行程序进行分析：

我们看到这与我们最初使用的 Python cProfiler 模块输出非常相似，我们用它来分析第一章中的 Mandelbrot 算法——只是现在，这专门告诉我们有关执行的所有 CUDA 操作。因此，当我们专门想要在 GPU 上进行优化时，我们可以使用它，而不必关心在主机上执行的任何 Python 或其他命令。（如果我们添加--print-gpu-trace命令行选项，我们可以进一步分析每个单独的 CUDA 内核操作，包括块和网格大小的启动参数。）

让我们再看一个技巧，帮助我们可视化程序所有操作的执行时间；我们将使用nvprof来转储一个文件，然后可以由 NVIDIA Visual Profiler 读取，以图形方式显示给我们。我们将使用上一章的示例multi-kernel_streams.py（在存储库的5下可用）来做这个。让我们回忆一下，这是我们对 CUDA 流概念的介绍示例之一，它允许我们同时执行和组织多个 GPU 操作。我们将使用-o命令行选项将输出转储到一个带有.nvvp文件后缀的文件中，如下所示：nvprof -o m.nvvp python multi-kernel_streams.py。现在我们可以使用nvvp m.nvvp命令将此文件加载到 NVIDIA Visual Profiler 中。

我们应该在所有 CUDA 流上看到一个时间线（记住，此程序中使用的内核名称为mult_ker）：

不仅可以看到所有的内核启动，还可以看到内存分配、内存复制和其他操作。这对于直观和视觉上理解程序如何随着时间使用 GPU 是有用的。

总结

我们在本章开始时看到了如何在 CUDA 内核中使用printf来输出来自各个线程的数据；我们特别看到了这对于调试代码有多么有用。然后，我们涵盖了 CUDA-C 中我们知识的一些空白，以便我们可以编写完整的测试程序，将其编译成适当的可执行二进制文件：这里有很多开销在我们之前是隐藏的，我们必须非常谨慎。接下来，我们看到了如何在 Nsight IDE 中创建和编译项目以及如何使用它进行调试。我们看到了如何在 CUDA 内核中停止我们设置的任何断点，并在不同的本地变量之间切换以查看各个线程。我们还使用了 Nsight 调试器来了解 warp 锁步属性以及为什么在 CUDA 内核中避免分支分歧很重要。最后，我们对 NVIDIA 命令行nvprof分析器和 Visual Profiler 进行了非常简要的概述，用于分析我们的 GPU 代码。

问题

1. 在我们编写的第一个 CUDA-C 程序中，我们在使用cudaMalloc在 GPU 上分配内存数组之后没有使用cudaDeviceSynchronize命令。为什么这是不必要的？（提示：回顾上一章。）

2. 假设我们有一个单个内核，它在由两个块组成的网格上启动，每个块有 32 个线程。假设第一个块中的所有线程都执行一个if语句，而第二个块中的所有线程都执行相应的else语句。第二个块中的所有线程是否必须像第一个块中的线程一样“锁步”执行if语句中的命令？

3. 如果我们执行类似的代码片段，只是在由 64 个线程执行的一个单一块组成的网格上执行，其中前 32 个线程执行一个if，而后 32 个执行一个else语句，会怎么样？

4. nvprof分析器可以为我们测量哪些 Python 的 cProfiler 无法测量的内容？

5. 列举一些我们可能更喜欢使用printf来调试 CUDA 内核的情境，以及其他一些情境，其中使用 Nsight 来调试 CUDA 内核可能更容易。

6. cudaSetDevice命令在 CUDA-C 中的目的是什么？

7. 为什么我们在每次 CUDA-C 中的内核启动或内存复制后都必须使用cudaDeviceSynchronize？

第七章：使用 Scikit-CUDA 与 CUDA 库

在本章中，我们将介绍三个用于简化数值和科学计算的标准 CUDA 库。我们将首先看一下cuBLAS，这是 NVIDIA 针对 CUDA 的基本线性代数子程序（BLAS）规范的实现。（cuBLAS 是 NVIDIA 对 BLAS 的各种优化的 CPU 实现的回应，例如免费/开源的 OpenBLAS 或英特尔的专有数学核心库。）接下来我们将看一下cuFFT，它可以在 GPU 上执行几乎每种快速傅里叶变换（FFT）的变体。我们将看看如何在图像处理中使用 cuFFT 进行滤波。然后我们将看一下cuSolver，它可以执行比 cuBLAS 中更复杂的线性代数运算，例如奇异值分解（SVD）或乔列斯基分解。

到目前为止，我们主要处理了一个作为我们与 CUDA 网关的单个 Python 模块——PyCUDA。虽然 PyCUDA 是一个非常强大和多功能的 Python 库，但它的主要目的是提供一个网关来编写、编译和启动 CUDA 内核，而不是提供一个接口给 CUDA 库。幸运的是，有一个免费的 Python 模块可用，它提供了一个用户友好的包装器接口给这些库。这就是 Scikit-CUDA。

虽然您不必了解 PyCUDA 甚至理解 GPU 编程就能欣赏 Scikit-CUDA，但它与 PyCUDA 兼容，例如，Scikit-CUDA 可以轻松地与 PyCUDA 的gpuarray类一起使用，这使您可以轻松地在我们自己的 CUDA 内核例程和 Scikit-CUDA 之间传递数据。此外，大多数例程也可以与 PyCUDA 的 stream 类一起使用，这将允许我们正确地同步我们自己的自定义 CUDA 内核和 Scikit-CUDA 的包装器。

请注意，除了这三个列出的库之外，Scikit-CUDA 还为专有的 CULA 库提供了包装器，以及开源的 MAGMA 库。这两者在功能上与官方的 NVIDIA 库有很多重叠。由于这些库在标准 CUDA 安装中默认未安装，我们将选择不在本章中涵盖它们。感兴趣的读者可以分别在www.culatools.com和icl.utk.edu/magma/了解更多关于 CULA 和 MAGMA 的信息。建议读者查看 Scikit-CUDA 的官方文档，网址为：media.readthedocs.org/pdf/scikit-cuda/latest/scikit-cuda.pdf。

本章的学习成果如下：

• 了解如何安装 Scikit-CUDA

• 了解标准 CUDA 库的基本目的和区别

• 了解如何使用低级 cuBLAS 函数进行基本线性代数

• 了解如何使用 SGEMM 和 DGEMM 操作来测量 GPU 在 FLOPS 中的性能

• 了解如何使用 cuFFT 在 GPU 上执行 1D 或 2D FFT 操作

• 了解如何使用 FFT 创建 2D 卷积滤波器，并将其应用于简单的图像处理

• 了解如何使用 cuSolver 执行奇异值分解（SVD）

• 了解如何使用 cuSolver 的 SVD 算法执行基本主成分分析

技术要求

本章需要一台安装了现代 NVIDIA GPU（2016 年及以后）的 Linux 或 Windows 10 PC，并安装了所有必要的 GPU 驱动程序和 CUDA Toolkit（9.0 及以后）。还需要一个合适的 Python 2.7 安装（如 Anaconda Python 2.7），其中包括 PyCUDA 模块。

本章的代码也可以在 GitHub 上找到，网址为：github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA。

有关先决条件的更多信息，请查看本书的前言。有关软件和硬件要求的更多信息，请查看github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA上的 README 文件。

安装 Scikit-CUDA

建议您直接从 GitHub 安装最新稳定版本的 Scikit-CUDA：github.com/lebedov/scikit-cuda。

将软件包解压缩到一个目录中，然后在此处打开命令行，并键入python setup.py install来安装模块。然后，您可以运行单元测试以确保已使用python setup.py test执行了正确的安装。（此方法建议 Windows 和 Linux 用户均使用。）或者，可以直接使用pip install scikit-cuda从 PyPI 存储库安装 Scikit-CUDA。

使用 cuBLAS 进行基本线性代数

我们将从学习如何使用 Scikit-CUDA 的 cuBLAS 包装器开始这一章。让我们花一点时间讨论 BLAS。BLAS（基本线性代数子程序）是一个基本线性代数库的规范，最早是在 1970 年代标准化的。BLAS 函数被分为几个类别，被称为级别。

Level 1 BLAS 函数包括纯粹在向量上的操作——向量-向量加法和缩放（也称为ax+y操作，或 AXPY），点积和范数。Level 2 BLAS 函数包括一般矩阵-向量操作（GEMV），例如矩阵与向量的乘法，而 Level 3 BLAS 函数包括“一般矩阵-矩阵”（GEMM）操作，例如矩阵-矩阵乘法。最初，这些库是在 1970 年代完全用 FORTRAN 编写的，因此您应该考虑到在使用和命名上可能存在一些看似过时的遗留问题，这可能对今天的新用户来说显得繁琐。

cuBLAS 是 NVIDIA 自己对 BLAS 规范的实现，当然是经过优化以充分利用 GPU 的并行性。Scikit-CUDA 提供了与 PyCUDA gpuarray对象兼容的 cuBLAS 包装器，以及与 PyCUDA 流兼容的包装器。这意味着我们可以通过 PyCUDA 将这些函数与我们自己的自定义 CUDA-C 内核耦合和接口，以及在多个流上同步这些操作。

使用 cuBLAS 进行 Level-1 AXPY

现在让我们从 cuBLAS 开始进行基本的 Level-1 ax + y（或 AXPY）操作。让我们停下来，回顾一下线性代数的一点，并思考这意味着什么。在这里，a被认为是一个标量；也就是说，一个实数，比如-10、0、1.345 或 100。x和y被认为是某个向量空间中的向量，。这意味着x和y是实数的 n 元组，因此在的情况下，这些值可以是[1,2,3]或[-0.345, 8.15, -15.867]。ax表示x的缩放乘以a，因此如果a是 10 且x是先前的第一个值，则ax是x的每个单独值乘以a；也就是[10, 20, 30]。最后，和ax + y表示我们将两个向量中每个槽的每个单独值相加以产生一个新的向量，假设y是给定的第二个向量，结果将如下所示-[9.655, 28.15, 14.133]。

现在让我们在 cuBLAS 中做这个。首先，让我们导入适当的模块：

import pycuda.autoinit
from pycuda import gpuarray
import numpy as np

现在让我们导入 cuBLAS：

from skcuda import cublas

我们现在可以设置我们的向量数组并将它们复制到 GPU。请注意，我们使用 32 位（单精度）浮点数：

a = np.float32(10)
x = np.float32([1,2,3])
y = np.float32([-.345,8.15,-15.867])
x_gpu = gpuarray.to_gpu(x)
y_gpu = gpuarray.to_gpu(y)

现在我们必须创建一个cuBLAS 上下文。这在本质上类似于 CUDA 上下文，在第五章中我们讨论过，只是这一次它被显式用于管理 cuBLAS 会话。cublasCreate函数创建一个 cuBLAS 上下文，并将其句柄作为输出。我们需要保存这个句柄，只要我们打算在此会话中使用 cuBLAS：

cublas_context_h = cublas.cublasCreate()

现在我们可以使用cublasSaxpy函数。S代表单精度，这是我们需要的，因为我们正在处理 32 位浮点数组：

cublas.cublasSaxpy(cublas_context_h, x_gpu.size, a, x_gpu.gpudata, 1, y_gpu.gpudata, 1)

让我们讨论一下我们刚刚做的事情。还要记住，这是一个直接包装到低级 C 函数的函数，因此输入可能更像 C 函数而不是真正的 Python 函数。简而言之，这执行了一个“AXPY”操作，最终将输出数据放入y_gpu数组中。让我们逐个讨论每个输入参数。

第一个输入始终是 CUDA 上下文句柄。然后我们必须指定向量的大小，因为这个函数最终将在 C 指针上操作；我们可以使用 gpuarray 的size参数来做到这一点。已经将我们的标量类型转换为 NumPy 的float32变量，我们可以将a变量直接作为标量参数传递。然后我们使用gpudata参数将x_gpu数组的底层 C 指针传递给这个函数。然后我们将第一个数组的步长设置为 1：步长指定每个输入值之间应该走多少步。 （相反，如果您正在使用来自行向矩阵的列的向量，则将步长设置为矩阵的宽度。）然后我们放入y_gpu数组的指针，并将其步长也设置为 1。

我们已经完成了计算；现在我们必须明确销毁我们的 cuBLAS 上下文：

cublas.cublasDestroy(cublas_context)

现在我们可以使用 NumPy 的allclose函数来验证这是否接近，就像这样：

print 'This is close to the NumPy approximation: %s' % np.allclose(a*x + y , y_gpu.get())

再次注意，最终输出被放入了y_gpu数组中，这也是一个输入。

始终记住，BLAS 和 CuBLAS 函数是就地操作，以节省时间和内存，而不是进行新的分配调用。这意味着输入数组也将用作输出！

我们刚刚看到如何使用cublasSaxpy函数执行AXPY操作。

让我们讨论突出的大写字母 S。正如我们之前提到的，这代表单精度，即 32 位实数浮点值（float32）。如果我们想要操作 64 位实数浮点值数组（NumPy 和 PyCUDA 中的float64），那么我们将使用cublasDaxpy函数；对于 64 位单精度复数值（complex64），我们将使用cublasCaxpy，而对于 128 位双精度复数值（complex128），我们将使用cublasZaxpy。

我们可以通过检查函数名称其余部分之前的字母来确定 BLAS 或 CuBLAS 函数操作的数据类型。使用单精度实数的函数总是以 S 开头，双精度实数以 D 开头，单精度复数以 C 开头，双精度复数以 Z 开头。

其他一级 cuBLAS 函数

让我们看看其他一些一级函数。我们不会深入介绍它们的操作，但步骤与我们刚刚介绍的类似：创建一个 cuBLAS 上下文，使用适当的数组指针调用函数（可以通过 PyCUDA 的gpuarray的gpudata参数访问），并相应地设置步长。另一件需要记住的事情是，如果函数的输出是单个值而不是数组（例如，点积函数），则函数将直接将该值输出到主机，而不是在必须从 GPU 中取出的内存数组中。（我们只会在这里介绍单精度实数版本，但其他数据类型的相应版本可以通过用适当的字母替换 S 来使用。）

我们可以对两个单精度实数gpuarray，v_gpu和w_gpu进行点积。再次，1 是为了确保我们在这个计算中使用步长 1！再次回想一下，点积是两个向量的逐点乘积的和：

dot_output = cublas.cublasSdot(cublas_context_h, v_gpu.size, v_gpu.gpudata, 1, w_gpu.gpudata, 1)

我们还可以执行向量的 L2 范数，就像这样（回想一下，对于一个向量x，这是它的 L2 范数，或长度，可以用公式来计算）：

l2_output = cublas.cublasSnrm2(cublas_context_h, v_gpu.size, v_gpu.gpudata, 1)

cuBLAS 中的二级 GEMV

让我们看看如何进行GEMV矩阵-向量乘法。对于一个m x n矩阵A，一个 n 维向量x，一个m维向量y，以及标量alpha和beta，这被定义为以下操作：

现在让我们在继续之前看一下函数的布局：

cublasSgemv(handle, trans, m, n, alpha, A, lda, x, incx, beta, y, incy)  

让我们逐个检查这些输入：

• handle指的是 cuBLAS 上下文句柄。

• trans指的是矩阵的结构——我们可以指定是否要使用原始矩阵、直接转置或共轭转置（对于复数矩阵）。这很重要要记住，因为这个函数将期望矩阵A以列主格式存储。

• m和n是我们想要使用的矩阵A的行数和列数。

• alpha是浮点值α。

• A是m x n矩阵A。

• lda指示矩阵的主维度，矩阵的总大小实际上是lda x n。这在列主格式中很重要，因为如果lda大于m，这可能会导致 cuBLAS 在尝试访问A的值时出现问题，因为该矩阵的基础结构是一个一维数组。

• 然后我们有x及其步长incx；x是被A相乘的向量的基础 C 指针。记住，x的大小必须是n；也就是说，A的列数。

• beta是浮点值β。

• 最后，我们有y及其步长incy作为最后的参数。我们应该记住，y的大小应该是m，或者A的行数。

让我们通过生成一个 10 x 100 的随机值矩阵A，和一个包含 100 个随机值的向量x来测试这个。我们将初始化y为一个包含 10 个零的矩阵。我们将 alpha 设置为 1，beta 设置为 0，以便直接进行矩阵乘法而不进行缩放：

m = 10
n = 100
alpha = 1
beta = 0
A = np.random.rand(m,n).astype('float32')
x = np.random.rand(n).astype('float32')
y = np.zeros(m).astype('float32')

我们现在必须将A转换为列主（或按列）格式。NumPy 默认将矩阵存储为行主（或按行）格式，这意味着用于存储矩阵的基础一维数组会遍历所有第一行的值，然后遍历所有第二行的值，依此类推。您应该记住，转置操作会将矩阵的列与其行交换。然而，结果将是转置矩阵的新一维数组将以列主格式表示原始矩阵。我们可以通过A.T.copy()这样做，将A的转置矩阵的副本以及x和y一起复制到 GPU 上：

A_columnwise = A.T.copy()
A_gpu = gpuarray.to_gpu(A_columnwise) 
x_gpu = gpuarray.to_gpu(x)
y_gpu = gpuarray.to_gpu(y)

由于我们现在已经正确地在 GPU 上存储了按列的矩阵，我们可以通过使用_CUBLAS_OP字典将trans变量设置为不进行转置：

trans = cublas._CUBLAS_OP['N']

由于矩阵的大小与我们想要使用的行数完全相同，我们现在将lda设置为m。x和y向量的步长再次为 1。我们现在已经设置好了所有需要的值，现在可以创建我们的 CuBLAS 上下文并存储它的句柄，就像这样：

lda = m 
incx = 1
incy = 1
handle = cublas.cublasCreate()

我们现在可以启动我们的函数。记住，A，x和y实际上是 PyCUDA gpuarray对象，所以我们必须使用gpudata参数输入到这个函数中。除了这个，这很简单：

cublas.cublasSgemv(handle, trans, m, n, alpha, A_gpu.gpudata, lda, x_gpu.gpudata, incx, beta, y_gpu.gpudata, incy)

现在我们可以销毁我们的 cuBLAS 上下文并检查返回值以确保它是正确的：

cublas.cublasDestroy(handle)
print 'cuBLAS returned the correct value: %s' % np.allclose(np.dot(A,x), y_gpu.get())

cuBLAS 中的三级 GEMM 用于测量 GPU 性能

现在我们将看看如何使用 CuBLAS 执行通用矩阵-矩阵乘法（GEMM）。实际上，我们将尝试制作一些比我们在 cuBLAS 中看到的最后几个示例更实用的东西-我们将使用这个作为我们的 GPU 性能指标，以确定它可以执行的每秒浮点运算次数（FLOPS）的数量，这将是两个单独的值：单精度和双精度的情况。使用 GEMM 是评估 FLOPS 中计算硬件性能的标准技术，因为它比使用纯时钟速度（MHz 或 GHz）更好地理解了纯计算能力。

如果您需要简要回顾，请回想一下我们在上一章中深入讨论了矩阵-矩阵乘法。如果您忘记了这是如何工作的，强烈建议您在继续本节之前复习一下这一章。

首先，让我们看看 GEMM 操作是如何定义的：

这意味着我们执行A和B的矩阵乘法，将结果缩放为alpha，然后加到我们已经通过beta缩放的C矩阵中，将最终结果放在C中。

让我们考虑一下执行实值 GEMM 操作的最终结果需要执行多少浮点运算，假设A是一个m x k（其中m是行，k是列）矩阵，B是一个k x n矩阵，C 是一个m x n矩阵。首先，让我们计算AB需要多少操作。让我们取A的一列并将其乘以B：这将导致m行中的每一行需要k次乘法和k-1次加法，这意味着这是m行上的km + (k-1)m总操作。B中有n列，因此计算AB将总共需要kmn + (k-1)mn = 2kmn - mn次操作。现在，我们使用alpha来缩放AB，这将是m**n次操作，因为这是矩阵AB的大小；类似地，通过beta缩放C是另外m**n次操作。最后，我们将这两个结果矩阵相加，这又是mn次操作。这意味着在给定的 GEMM 操作中，我们将有2kmn - mn + 3mn = 2kmn + 2mn = 2mn(k+1)次浮点运算。

现在我们唯一需要做的就是运行一个计时的 GEMM 操作，注意矩阵的不同大小，并将2kmn + 2mn除以总时间持续时间来计算我们 GPU 的 FLOPS。结果数字将非常大，因此我们将以 GFLOPS 的形式表示这一点-也就是说，每秒可以计算多少十亿（10⁹）次操作。我们可以通过将 FLOPS 值乘以 10^(-9)来计算这一点。

现在我们准备开始编写代码。让我们从我们的导入语句开始，以及time函数：

import pycuda.autoinit
from pycuda import gpuarray
import numpy as np
from skcuda import cublas
from time import time

现在我们将为我们的矩阵大小设置m，n和k变量。我们希望我们的矩阵相对较大，以便时间持续足够长，以避免除以 0 的错误。以下值对于截至 2018 年中旬或更早发布的任何 GPU 来说应该足够了；拥有更新卡的用户可能考虑增加这些值：

m = 5000
n = 10000
k = 10000

现在我们将编写一个计算单精度和双精度 GFLOPS 的函数。如果我们希望使用双精度，则将输入值设置为'D'，否则设置为'S'：

def compute_gflops(precision='S'):

if precision=='S':
    float_type = 'float32'
elif precision=='D':
    float_type = 'float64'
else:
    return -1

现在让我们生成一些随机矩阵，这些矩阵具有我们将用于计时的适当精度。GEMM 操作与我们之前看到的 GEMV 操作类似，因此我们必须在将它们复制到 GPU 之前对其进行转置。（由于我们只是在计时，这一步并不是必要的，但记住这一点是个好习惯。）

我们将为 GEMM 设置一些其他必要的变量，这些变量在这一点上应该是不言自明的（transa，lda，ldb等）：

A = np.random.randn(m, k).astype(float_type)
B = np.random.randn(k, n).astype(float_type)
C = np.random.randn(m, n).astype(float_type)
A_cm = A.T.copy()
B_cm = B.T.copy()
C_cm = C.T.copy()
A_gpu = gpuarray.to_gpu(A_cm)
B_gpu = gpuarray.to_gpu(B_cm)
C_gpu = gpuarray.to_gpu(C_cm)
alpha = np.random.randn()
beta = np.random.randn()
transa = cublas._CUBLAS_OP['N']
transb = cublas._CUBLAS_OP['N']
lda = m
ldb = k
ldc = m

现在我们可以开始计时了！首先，我们将创建一个 cuBLAS 上下文：

t = time()
handle = cublas.cublasCreate()

现在我们将启动 GEMM。请记住，对于实数情况有两个版本：cublasSgemm用于单精度，cublasDgemm用于双精度。我们可以使用一个小小的 Python 技巧执行适当的函数：我们将用cublas%sgemm和适当的参数写一个字符串，然后通过附加% precision将%s替换为 D 或 S。然后我们将使用exec函数将这个字符串作为 Python 代码执行，就像这样：

exec('cublas.cublas%sgemm(handle, transa, transb, m, n, k, alpha, A_gpu.gpudata, lda, B_gpu.gpudata, ldb, beta, C_gpu.gpudata, ldc)' % precision)

现在我们可以销毁 cuBLAS 上下文，并得到我们计算的最终时间：

cublas.cublasDestroy(handle)
t = time() - t

然后我们需要使用我们推导出的方程计算 GFLOPS，并将其作为这个函数的输出返回：

gflops = 2*m*n*(k+1)*(10**-9) / t 
return gflops

现在我们可以设置我们的主函数。我们将输出单精度和双精度情况下的 GFLOPS：

if __name__ == '__main__':
    print 'Single-precision performance: %s GFLOPS' % compute_gflops('S')
    print 'Double-precision performance: %s GFLOPS' % compute_gflops('D')

现在在运行这个程序之前，让我们做一点功课——去www.techpowerup.com搜索你的 GPU，然后注意两件事——单精度浮点性能和双精度浮点性能。我现在使用的是 GTX 1050，它的列表声称在单精度上有 1,862 GFLOPS 性能，在双精度上有 58.20 GFLOPS 性能。让我们现在运行这个程序，看看这是否符合事实：

看哪，它成功了！

这个程序也可以在这本书的存储库目录下的cublas_gemm_flops.py文件中找到。

使用 cuFFT 进行快速傅里叶变换

现在让我们看看如何使用 cuFFT 进行一些基本的快速傅里叶变换（FFT）。首先，让我们简要回顾一下傅里叶变换到底是什么。如果你上过高级微积分或分析课程，你可能已经见过傅里叶变换被定义为一个积分公式，就像这样：

这个程序将f作为x的时间域函数。这给了我们一个相应的频域函数，对应于"ξ"。这事实上是一个极其有用的工具，几乎触及到所有科学和工程的分支。

让我们记住积分可以被看作是一个求和；同样地，有一个对应的离散、有限版本的傅里叶变换，称为离散傅里叶变换（DFT）。这对长度为n的向量进行操作，并允许它们在频域中进行分析或修改。x的 DFT 被定义如下：

换句话说，我们可以将一个向量x乘以复杂的N x N矩阵 

（这里，k对应行号，n对应列号）来找到它的 DFT。我们还应该注意到逆公式，它让我们从它的 DFT 中检索x（在这里用x的 DFT 替换y，输出将是原始的x）：

通常，计算矩阵-向量操作的计算复杂度对于长度为N的向量是 O(N²)。然而，由于 DFT 矩阵中的对称性，这总是可以通过使用 FFT 减少到 O(N log N)。让我们看看如何使用 FFT 与 CuBLAS，然后我们将继续一个更有趣的例子。

一个简单的一维 FFT

让我们首先看看如何使用 cuBLAS 计算简单的一维 FFT。首先，我们将简要讨论 Scikit-CUDA 中的 cuFFT 接口。

这里有两个子模块，我们可以使用cufft和fft访问 cuFFT 库。cufft包括了一系列 cuFFT 库的低级封装，而fft提供了一个更加用户友好的接口；在本章中我们将只使用fft。

让我们从适当的导入开始，记得包括 Scikit-CUDA 的fft子模块：

import pycuda.autoinit
from pycuda import gpuarray
import numpy as np
from skcuda import fft

现在我们将设置一些随机数组并将其复制到 GPU。我们还将设置一个空的 GPU 数组，用于存储 FFT（请注意，我们使用实数 float32 数组作为输入，但输出将是 complex64 数组，因为傅立叶变换总是复值）：

x = np.asarray(np.random.rand(1000), dtype=np.float32 )
x_gpu = gpuarray.to_gpu(x)
x_hat = gpuarray.empty_like(x_gpu, dtype=np.complex64)

我们现在将为正向 FFT 变换设置一个 cuFFT 计划。这是 cuFFT 用来确定变换的形状，以及输入和输出数据类型的对象：

plan = fft.Plan(x_gpu.shape,np.float32,np.complex64)

我们还将为逆 FFT 计划对象设置一个计划。请注意，这一次我们从complex64到实数float32：

inverse_plan = fft.Plan(x.shape, in_dtype=np.complex64, out_dtype=np.float32)

现在，我们必须将x_gpu的正向 FFT 转换为x_hat，并将x_hat的逆 FFT 转换回x_gpu。请注意，在逆 FFT 中设置了scale=True；我们这样做是为了告诉 cuFFT 将逆 FFT 按 1/N 进行缩放：

fft.fft(x_gpu, x_hat, plan)
fft.ifft(x_hat, x_gpu, inverse_plan, scale=True)

我们现在将检查x_hat与x的 NumPy FFT，以及x_gpu与x本身：

y = np.fft.fft(x)
print 'cuFFT matches NumPy FFT: %s' % np.allclose(x_hat.get(), y, atol=1e-6)
print 'cuFFT inverse matches original: %s' % np.allclose(x_gpu.get(), x, atol=1e-6)

如果你运行这个程序，你会发现x_hat与y不匹配，然而，莫名其妙的是，x_gpu与x匹配。这是怎么可能的？好吧，让我们记住x是实数；如果你看一下离散傅立叶变换是如何计算的，你可以数学上证明实向量的输出在 N/2 之后会重复为它们的复共轭。虽然 NumPy FFT 会完全计算这些值，但 cuFFT 通过只计算输入为实数时的输出的前半部分来节省时间，并将其余输出设置为0。你应该通过检查前面的变量来验证这一点。

因此，如果我们将前面代码中的第一个打印语句更改为仅比较 CuFFT 和 NumPy 之间的前 N/2 个输出，那么这将返回 true：

print 'cuFFT matches NumPy FFT: %s' % np.allclose(x_hat.get()[0:N//2], y[0:N//2], atol=1e-6)

使用 FFT 进行卷积

我们现在将看看如何使用 FFT 执行卷积。首先让我们回顾一下卷积的确切定义：给定两个一维向量x和y，它们的卷积定义如下：

这对我们很有意义，因为如果x是一些长的连续信号，而y只有一小部分局部非零值，那么y将作为x的滤波器；这本身有许多应用。首先，我们可以使用滤波器平滑信号x（在数字信号处理和图像处理中很常见）。我们还可以使用它来收集信号x的样本，以表示信号或压缩它（在数据压缩或压缩感知领域很常见），或者使用滤波器为机器学习中的信号或图像识别收集特征。这个想法构成了卷积神经网络的基础。

当然，计算机无法处理无限长的向量（至少目前还不能），所以我们将考虑循环卷积。在循环卷积中，我们处理两个长度为n的向量，其索引小于 0 或大于 n-1 的部分将会环绕到另一端；也就是说，x[-1] = x[n-1]，x[-2] = x[n-2]，x[n] = x[0]，x[n+1] = x[1]，依此类推。我们定义x和y的循环卷积如下：

事实证明，我们可以很容易地使用 FFT 执行循环卷积；我们可以对x和y执行 FFT，对输出进行逐点乘法，然后对最终结果执行逆 FFT。这个结果被称为卷积定理，也可以表示如下：

我们将在两个维度上进行这个操作，因为我们希望将结果应用到信号处理上。虽然我们只看到了一维卷积和 FFT 的数学运算，但二维卷积和 FFT 的工作方式与一维类似，只是索引更复杂一些。然而，我们选择跳过这一点，以便我们可以直接进入应用。

使用 cuFFT 进行 2D 卷积

现在我们将制作一个小程序，使用基于 cuFFT 的二维卷积对图像进行高斯滤波。高斯滤波是一种使用所谓的高斯滤波器平滑粗糙图像的操作。之所以这样命名，是因为它基于统计学中的高斯（正态）分布。这是高斯滤波器在两个维度上以标准偏差σ定义的方式：

当我们用滤波器对离散图像进行卷积时，有时我们会将滤波器称为卷积核。通常，图像处理工程师会简单地称之为核，但由于我们不想将其与 CUDA 核混淆，我们将始终使用完整术语卷积核。在这里，我们将使用离散版本的高斯滤波器作为我们的卷积核。

让我们从适当的导入开始；请注意，我们将在这里使用 Scikit-CUDA 子模块linalg。这将为我们提供比 cuBLAS 更高级的接口。由于我们在这里处理图像，我们还将导入 Matplotlib 的pyplot子模块。还要注意，我们将在此处使用 Python 3 风格的除法，从第一行开始；这意味着如果我们使用/运算符除两个整数，那么返回值将是浮点数，无需类型转换（我们使用//运算符进行整数除法）：

from __future__ import division
import pycuda.autoinit
from pycuda import gpuarray
import numpy as np
from skcuda import fft
from skcuda import linalg
from matplotlib import pyplot as plt

让我们立即开始编写卷积函数。这将接受两个相同大小的 NumPy 数组x和y。我们将把它们转换为 complex64 数组，然后如果它们的大小不相同，就返回-1：

def cufft_conv(x , y):
    x = x.astype(np.complex64)
    y = y.astype(np.complex64)

    if (x.shape != y.shape):
        return -1

现在我们将设置我们的 FFT 计划和逆 FFT 计划对象：

plan = fft.Plan(x.shape, np.complex64, np.complex64)
inverse_plan = fft.Plan(x.shape, np.complex64, np.complex64)

现在我们可以将我们的数组复制到 GPU。我们还将设置一些空数组，用于保存这些数组的 FFT 的适当大小，另外一个数组将保存最终卷积的输出，out_gpu：

 x_gpu = gpuarray.to_gpu(x)
 y_gpu = gpuarray.to_gpu(y)

 x_fft = gpuarray.empty_like(x_gpu, dtype=np.complex64)
 y_fft = gpuarray.empty_like(y_gpu, dtype=np.complex64)
 out_gpu = gpuarray.empty_like(x_gpu, dtype=np.complex64)

现在我们可以进行我们的 FFT：

fft.fft(x_gpu, x_fft, plan)
fft.fft(y_gpu, y_fft, plan)

我们现在将使用linalg.multiply函数在x_fft和y_fft之间执行逐点（Hadamard）乘法。我们将设置overwrite=True，以便将最终值写入y_fft：

linalg.multiply(x_fft, y_fft, overwrite=True)

现在我们将调用逆 FFT，将最终结果输出到out_gpu。我们将这个值传输到主机并返回它：

fft.ifft(y_fft, out_gpu, inverse_plan, scale=True)
conv_out = out_gpu.get()
return conv_out

我们还没有完成。我们的卷积核将比输入图像小得多，因此我们将不得不调整我们的两个 2D 数组的大小（卷积核和图像），使它们相等，并在它们之间执行逐点乘法。我们不仅应该确保它们相等，还需要确保我们在数组上执行零填充，并适当地将卷积核居中。零填充意味着我们在图像的两侧添加零缓冲区，以防止环绕错误。如果我们使用 FFT 来执行卷积，记住它是循环卷积，因此边缘将始终环绕。当我们完成卷积后，我们可以去除图像外部的缓冲区，得到最终的输出图像。

让我们创建一个名为conv_2d的新函数，它接受卷积核ker和图像img。填充后的图像大小将是（2*ker.shape[0] + img.shape[0]，2*ker.shape[1] + img.shape[1]）。让我们首先设置填充的卷积核。我们将创建一个这个大小的零矩阵，然后将左上角子矩阵设置为我们的卷积核，如下所示：

def conv_2d(ker, img):

    padded_ker = np.zeros( (img.shape[0] + 2*ker.shape[0], img.shape[1] + 2*ker.shape[1] )).astype(np.float32)
    padded_ker[:ker.shape[0], :ker.shape[1]] = ker

现在我们需要移动卷积核，使其中心精确地位于坐标（0,0）处。我们可以使用 NumPy 的roll命令来实现这一点：

padded_ker = np.roll(padded_ker, shift=-ker.shape[0]//2, axis=0)
padded_ker = np.roll(padded_ker, shift=-ker.shape[1]//2, axis=1)

现在我们需要对输入图像进行填充：

padded_img = np.zeros_like(padded_ker).astype(np.float32)
padded_img[ker.shape[0]:-ker.shape[0], ker.shape[1]:-ker.shape[1]] = img

现在我们有两个大小相同且格式正确的数组。我们现在可以使用我们刚刚在这里编写的cufft_conv函数：

out_ = cufft_conv(padded_ker, padded_img)

现在我们可以去除图像外部的零缓冲区。然后返回结果：

output = out_[ker.shape[0]:-ker.shape[0], ker.shape[1]:-ker.shape[1]]

return output

我们还没有完成。让我们编写一些小函数来设置我们的高斯滤波器，然后我们可以继续将其应用于图像。我们可以使用 lambda 函数一行代码来编写基本滤波器本身：

gaussian_filter = lambda x, y, sigma : (1 / np.sqrt(2*np.pi*(sigma**2)) )*np.exp( -(x**2 + y**2) / (2 * (sigma**2) ))

现在我们可以编写一个使用此滤波器输出离散卷积核的函数。卷积核的高度和长度将为2*sigma + 1，这是相当标准的：

请注意，我们通过将其值求和为total_并将其除以来规范化高斯核的值。

def gaussian_ker(sigma):
    ker_ = np.zeros((2*sigma+1, 2*sigma+1))
    for i in range(2*sigma + 1):
        for j in range(2*sigma + 1):
            ker_[i,j] = gaussian_filter(i - sigma, j - sigma, sigma)
    total_ = np.sum(ker_.ravel())
    ker_ = ker_ */* total*_* return ker_

我们现在准备在图像上测试这个！作为我们的测试案例，我们将使用高斯滤波来模糊这本书编辑Akshada Iyer的彩色 JPEG 图像。（此图像在 GitHub 存储库的Chapter07目录中，文件名为akshada.jpg。）我们将使用 Matplotlib 的imread函数来读取图像；默认情况下，这将存储为 0 到 255 范围内的无符号 8 位整数数组。我们将将其强制转换为浮点数组并对其进行规范化，以便所有值的范围为 0 到 1。

对于本书印刷版的读者注意：尽管本书的印刷版是灰度图像，但这是一幅彩色图像。然后，我们将设置一个空的零数组，用于存储模糊的图像：

if __name__ == '__main__':
    akshada = np.float32(plt.imread('akshada.jpg')) / 255
    akshada_blurred = np.zeros_like(akshada)

让我们设置我们的卷积核。在这里，标准差为 15 应该足够：

ker = gaussian_ker(15)

现在我们可以模糊图像。由于这是一幅彩色图像，我们将不得不分别对每个颜色层（红色、绿色和蓝色）应用高斯滤波；这在图像数组的第三维中进行索引：

for k in range(3):
    akshada_blurred[:,:,k] = conv_2d(ker, akshada[:,:,k])

现在让我们通过使用一些 Matplotlib 技巧并排查看 Before 和 After 图像：

fig, (ax0, ax1) = plt.subplots(1,2)
fig.suptitle('Gaussian Filtering', fontsize=20)
ax0.set_title('Before')
ax0.axis('off')
ax0.imshow(akshada)
ax1.set_title('After')
ax1.axis('off')
ax1.imshow(akshada_blurred)
plt.tight_layout()
plt.subplots_adjust(top=.85)
plt.show()

现在我们可以运行程序并观察高斯滤波的效果：

该程序在此书的存储库中的Chapter07目录中的名为conv_2d.py的文件中可用。

使用 Scikit-CUDA 中的 cuSolver

我们现在将看看如何使用 Scikit-CUDA 的linalg子模块中的 cuSolver。同样，这为 cuBLAS 和 cuSolver 提供了一个高级接口，因此我们不必陷入细节。

正如我们在介绍中指出的，cuSolver 是一个用于执行比 cuBLAS 更高级的线性代数运算的库，例如奇异值分解、LU/QR/Cholesky 分解和特征值计算。由于 cuSolver、cuBLAS 和 cuFFT 一样，是另一个庞大的库，我们只会花时间来看数据科学和机器学习中最基本的操作之一——SVD。

如果您想进一步了解此库，请参考 NVIDIA 关于 cuSOLVER 的官方文档：docs.NVIDIA.com/cuda/cusolver/index.html。

奇异值分解（SVD）

SVD 接受任何m x n矩阵A，然后返回三个矩阵—U、Σ和V。在这里，U是一个m x m酉矩阵，Σ是一个m x n对角矩阵，V是一个n x n酉矩阵。这里，酉表示矩阵的列形成一个正交归一基；对角表示矩阵中的所有值都为零，除了可能沿着对角线的值。

SVD 的重要性在于它将A分解为这些矩阵，使得我们有A = UΣV^T；此外，Σ对角线上的值将全部为正值或零，并且被称为奇异值。我们很快将看到一些应用，但您应该记住，SVD 的计算复杂度为 O(mn²)——对于大矩阵来说，使用 GPU 绝对是一个好主意，因为这个算法是可并行化的。

现在让我们看看如何计算矩阵的 SVD。让我们进行适当的导入语句：

import pycuda.autoinit
from pycuda import gpuarray
import numpy as np
from skcuda import linalg

现在我们将生成一个相对较大的随机矩阵并将其传输到 GPU：

a = np.random.rand(1000,5000).astype(np.float32)
a_gpu = gpuarray.to_gpu(a)

现在我们可以执行 SVD。这将有三个输出，对应于我们刚刚描述的矩阵。第一个参数将是我们刚刚复制到 GPU 的矩阵数组。然后我们需要指定我们要使用 cuSolver 作为此操作的后端：

U_d, s_d, V_d = linalg.svd(a_gpu,  lib='cusolver')

现在让我们将这些数组从 GPU 复制到主机：

U = U_d.get()
s = s_d.get()
V = V_d.get()

s实际上存储为一维数组；我们将不得不创建一个大小为 1000 x 5000 的零矩阵，并将这些值沿对角线复制。我们可以使用 NumPy 的diag函数来做到这一点，再加上一些数组切片：

S = np.zeros((1000,5000))
S[:1000,:1000] = np.diag(s)

我们现在可以在主机上使用 NumPy 的dot函数对这些值进行矩阵相乘，以验证它们是否与我们的原始数组匹配：

print 'Can we reconstruct a from its SVD decomposition? : %s' % np.allclose(a, np.dot(U, np.dot(S, V)), atol=1e-5)

由于我们只使用 float32，并且我们的矩阵相对较大，引入了一些数值误差；我们不得不将“容差”级别（atol）设置得比平常高一点，但仍然足够小，以验证这两个数组是足够接近的。

使用 SVD 进行主成分分析（PCA）

主成分分析（PCA）是主要用于降维的工具。我们可以使用它来查看数据集，并找出哪些维度和线性子空间最显著。虽然有几种实现方法，但我们将向您展示如何使用 SVD 执行 PCA。

我们将这样做——我们将使用一个存在于 10 个维度中的数据集。我们将首先创建两个在前面有很大权重的向量，其他位置为 0：

vals = [ np.float32([10,0,0,0,0,0,0,0,0,0]) , np.float32([0,10,0,0,0,0,0,0,0,0]) ]

然后我们将添加 9000 个额外的向量：其中 6000 个将与前两个向量相同，只是加了一点随机白噪声，剩下的 3000 个将只是随机白噪声：

for i in range(3000):
    vals.append(vals[0] + 0.001*np.random.randn(10))
    vals.append(vals[1] + 0.001*np.random.randn(10))
    vals.append(0.001*np.random.randn(10))

我们现在将vals列表转换为float32的 NumPy 数组。我们对行进行平均，并从每行中减去这个值。（这是 PCA 的必要步骤。）然后我们转置这个矩阵，因为 cuSolver 要求输入矩阵的行数少于或等于列数：

vals = np.float32(vals)
vals = vals - np.mean(vals, axis=0)
v_gpu = gpuarray.to_gpu(vals.T.copy())

我们现在将运行 cuSolver，就像我们之前做的那样，并将输出值从 GPU 复制出来：

U_d, s_d, V_d = linalg.svd(v_gpu, lib='cusolver')

u = U_d.get()
s = s_d.get()
v = V_d.get()

现在我们准备开始我们的调查工作。让我们打开 IPython，仔细查看u和s。首先，让我们看看 s；它的值实际上是主要值的平方根，所以我们将它们平方然后看一看：

您会注意到，前两个主要值的数量级为 10⁵，而其余的分量的数量级为 10^(-3)。这告诉我们，实际上只有一个二维子空间与这些数据相关，这并不令人意外。这些是第一个和第二个值，它们将对应于第一个和第二个主成分，即相应的向量。让我们来看看这些向量，它们将存储在U中：

！

您会注意到这两个向量在前两个条目中有很大的权重，数量级为 10^(-1)；其余的条目都是 10^(-6)或更低，相对不相关。考虑到我们在前两个条目中使数据偏向，这正是我们应该期望的。这就是 PCA 背后的思想。

摘要

我们从查看如何使用 Scikit-CUDA 库的 cuBLAS 包装器开始了本章；在这里，我们必须记住许多细节，比如何使用列主存储，或者输入数组是否会被就地覆盖。然后，我们看了如何使用 Scikit-CUDA 的 cuFFT 执行一维和二维 FFT，以及如何创建一个简单的卷积滤波器。然后，我们向您展示了如何将其应用于图像的简单高斯模糊效果。最后，我们看了如何使用 cuSolver 在 GPU 上执行奇异值分解（SVD），这通常是一个非常计算密集的操作，但在 GPU 上可以很好地并行化。我们通过查看如何使用 SVD 进行基本的 PCA 来结束本章。

问题

1. 假设你得到了一个工作，需要将一些旧的遗留 FORTRAN BLAS 代码转换成 CUDA。你打开一个文件，看到一个名为 SBLAH 的函数，另一个名为 ZBLEH。你能在不查找的情况下告诉这两个函数使用的数据类型吗？

2. 你能修改 cuBLAS level-2 GEMV 示例，直接将矩阵A复制到 GPU，而不是在主机上进行转置以设置为列优先吗？

3. 使用 cuBLAS 32 位实数点积（cublasSdot）来实现使用一个按行矩阵和一个步幅为 1 的向量进行矩阵-向量乘法。

4. 使用cublasSdot实现矩阵-矩阵乘法。

5. 你能实现一种精确测量性能测量示例中的 GEMM 操作的方法吗？

6. 在一维 FFT 的示例中，尝试将x强制转换为complex64数组，然后将 FFT 和逆 FFT 计划都设置为complex64值。然后确认np.allclose(x, x_gpu.get())是否为真，而不检查数组的前一半。你认为为什么现在这样做会有效？

7. 请注意，在卷积示例中，模糊图像周围有一个暗边。为什么模糊图像中有这个暗边，而原始图像中没有呢？你能想到一种方法来减轻这个问题吗？

第八章：CUDA 设备函数库和 Thrust

在上一章中，我们对通过 Scikit-CUDA 包装器模块在 CUDA 中可用的库进行了相当广泛的概述。我们现在将看一下另外一些库，我们将不得不直接从 CUDA C 中使用这些库，而不是像 Scikit-CUDA 中的包装器那样。我们将首先看一下两个标准库，其中包含我们可以从任何 CUDA C 内核中调用的设备函数 cuRAND 和 CUDA Math API。通过学习如何使用这些库，我们将了解如何在蒙特卡罗积分的上下文中使用这些库。蒙特卡罗积分是一种众所周知的随机方法，可以提供来自微积分的定积分值的估计。我们首先将看一个基本示例，演示如何使用 cuRAND 实现简单的蒙特卡罗方法来对π的值进行基本估计（如众所周知的常数，π=3.14159...），然后我们将着手进行一个更有雄心的项目，我们将构建一个 Python 类，可以对任意数学函数执行定积分，并使用 Math API 创建这样的函数。我们还将看一下如何在设计这个类时有效地使用元编程的一些思想。

然后，我们将再次使用 Thrust C++库来编写一些纯 CUDA 程序。Thrust 是一个提供 C++模板容器的库，类似于 C++标准模板库（STL）中的容器。这将使我们能够以更接近 PyCUDA 的gpuarray和 STL 的向量容器的更自然的方式从 C++中操作 CUDA C 数组。这将使我们免受在 CUDA C 中以前不断使用指针（如mallocs和frees）的困扰。

在本章中，我们将讨论以下主题：

• 理解种子在生成伪随机数列表中的作用

• 在 CUDA 内核中使用 cuRAND 设备函数生成随机数

• 理解蒙特卡罗积分的概念

• 在 Python 中使用基于字典的字符串格式化进行元编程

• 使用 CUDA Math API 设备函数库

• 理解 functor 是什么

• 在纯 CUDA C 编程时使用 Thrust 向量容器

技术要求

本章需要具有现代 NVIDIA GPU（2016 年至今）的 Linux 或 Windows 10 PC，并安装了所有必要的 GPU 驱动程序和 CUDA Toolkit（9.0 及以上）。还需要适当的 Python 2.7 安装（如 Anaconda Python 2.7），并安装了 PyCUDA 模块。

本章的代码也可以在 GitHub 上找到，网址为github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA.

有关本章先决条件的更多信息，请查看本书的前言。有关软件和硬件要求，请查看github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA上的 README。

cuRAND 设备函数库

让我们从 cuRAND 开始。这是一个标准的 CUDA 库，用于在 CUDA 内核中按线程生成伪随机值，通过调用每个线程内核中的设备函数进行初始化和调用。让我们再次强调，这是一个伪随机值序列——因为数字硬件始终是确定性的，从不是随机或任意的，我们使用算法从初始种子值生成一系列表面上随机的值。通常，我们可以将种子值设置为真正的随机值（例如毫秒级的时钟时间），这将产生一系列随机值。这些生成的随机值与由相同种子生成的序列中的先前或未来值没有相关性，尽管当您组合从不同种子生成的值时，可能会出现相关性和重复。因此，您必须小心，希望彼此随机的值是由相同的种子生成的。

让我们从curand_init的函数原型开始，我们将使用适当的种子进行初始化：

__device__ void curand_init ( unsigned long long seed, unsigned long long sequence, unsigned long long offset, curandState_t *state)

在这里，所有的输入都是无符号长整型，在 C 中是无符号（非负值）的 64 位整数。首先，我们可以看到seed，这当然是种子值。一般来说，您将使用时钟值或某种变化来设置这个值。然后我们看到一个称为sequence的值，正如我们之前所述，cuRAND 生成的值只有在它们由相同的种子值生成时才是真正的数学上相互随机的。因此，如果我们有多个线程使用相同的种子值，我们使用sequence来指示当前线程使用长度为 2¹⁹⁰的随机数子序列的哪个子序列，而我们使用offset来指示从这个子序列的哪个点开始；这将在每个线程中生成所有数学上相互随机且没有相关性的值。最后，最后一个参数是指向curandState_t对象的指针；这将跟踪我们在伪随机数序列中的位置。

初始化类对象后，您将通过调用适当的设备函数从适当的随机分布生成随机值。最常见的两种分布是均匀分布和正态（高斯）分布。均匀分布（curand_uniform，在 cuRAND 中）是一个输出值在给定范围内都是等概率的函数：也就是说，对于 0 到 1 的均匀分布，值在 0 到 0.1 之间的概率是 10%，或者在 0.9 到 1 之间，或者在任何两个相距 0.1 的点之间。正态分布（curand_normal，在 cuRAND 中）具有以特定均值为中心的值，这些值将根据分布的标准差分布在众所周知的钟形曲线上。 （curand_normal的默认均值为0，标准差为 1，在 cuRAND 中，因此必须手动移位和缩放为其他值。）cuRAND 支持的另一个众所周知的分布是泊松分布（curand_poisson），用于对随机事件的发生进行建模。

在接下来的部分中，我们将主要研究如何在均匀分布的背景下使用 cuRAND，因为它们适用于蒙特卡罗积分。鼓励有兴趣学习如何使用 cuRAND 更多功能的读者查看 NVIDIA 的官方文档。

用蒙特卡洛法估算π

首先，我们将运用我们对 cuRAND 的新知识来估算众所周知的数学常数π，或圆周率，这当然是永不停止的无理数 3.14159265358979...

然而，要得到一个估计值，我们需要花点时间思考这意味着什么。让我们想想一个圆。记住，圆的半径是从圆心到圆上任意一点的长度；通常用 R 表示。直径被定义为 D = 2R，周长 C 是围绕圆的长度。然后，π 被定义为 π = C / D。我们可以使用欧几里得几何来找到圆的面积公式，结果是 A = πR²。现在，让我们考虑一个半径为 R 的圆被内切在边长为 2R 的正方形中：

因此，当然，我们知道正方形的面积是 (2R)² = 4R²。让我们考虑 R=1，这样我们就知道圆的面积恰好是 π，而正方形的面积恰好是 4。让我们进一步假设并声明，圆和正方形都以 (0,0) 为中心。现在，让我们在正方形内取一个完全随机的值 (x,y)，并查看它是否落在圆内。我们如何做到这一点？通过应用勾股定理公式：我们通过检查 x² + y² 是否小于或等于 1 来做到这一点。让我们用 iters 表示我们选择的随机点的总数，用 hits 表示命中的次数。

让我们再多想一下：在圆内选择一个点的概率应该与圆的面积与矩形的面积之比成比例；在这里，这是 π / 4。然而，如果我们选择一个非常大的随机点值，注意到我们将得到以下近似值：

这正是我们将估计 π 的方法！在我们能得出合理的 π 估计之前，我们将不得不进行非常多的迭代，但请注意这是多么好的可并行化：我们可以在不同的线程中检查“命中”，将总迭代次数分配给不同的线程。在一天结束时，我们只需将所有线程中的命中总数相加，即可得到我们的估计值。

现在，我们可以开始编写一个程序来进行蒙特卡洛估计。让我们首先导入我们在 PyCUDA 程序中需要的常规 Python 模块，再加上 SymPy 中的一个模块：

SymPy 用于在 Python 中进行完美的 符号 计算，因此当我们有非常大的整数时，我们可以使用 Rational 函数来对除法进行更准确的浮点估计。

import pycuda.autoinit
import pycuda.driver as drv
from pycuda import gpuarray
from pycuda.compiler import SourceModule
import numpy as np
from sympy import Rational

现在，当我们构建内核时，我们必须做一些与正常情况不同的事情：我们需要在 SourceModule 中设置选项 no_extern_c=True。这会修改代码的编译方式，以便我们的代码可以正确地与 cuRAND 库所需的 C++ 代码链接。然后我们开始编写我们的内核并包含适当的头文件：

ker = SourceModule(no_extern_c=True, source='''
#include <curand_kernel.h>

现在，让我们包含一个用于勾股定理距离的宏。由于我们只是检查这个值是否等于或小于 1，因此可以省略平方根。我们将使用大量的无符号 64 位整数，因此让我们再定义一个宏，以免我们一遍又一遍地输入 unsigned long long：

#define _PYTHAG(a,b) (a*a + b*b)
#define ULL unsigned long long

现在，我们可以设置我们的内核。根据 PyCUDA 的性质，这将不得不编译为接口的真正的 C 函数，而不是 C++ 函数。我们可以通过 extern "C" 块来实现这一点：

extern "C" {

现在，我们可以定义我们的内核。我们将有两个参数：一个是 iters，它是每个线程的总迭代次数，另一个是一个数组，将保存每个线程的命中总数。我们将需要一个 curandState 对象：

__global__ void estimate_pi(ULL iters, ULL * hits)
{
    curandState cr_state;

让我们用一个名为 tid 的整数来保存全局线程 ID：

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

clock()是一个设备函数，输出当前时间到毫秒。我们可以将tid添加到clock()的输出中，以获得每个线程的唯一种子。我们不需要使用不同的子序列或偏移量，所以让我们都设置为 0。我们还将在这里仔细地将所有内容强制转换为 64 位无符号整数：

curand_init( (ULL) clock() + (ULL) tid, (ULL) 0, (ULL) 0, &cr_state);

让我们设置x和y值以保存矩形中的随机点：

float x, y;

然后我们将迭代iters次，看看我们在圆中得到了多少次命中。我们使用curand_uniform(&cr_state)生成这些。请注意，我们可以在 0 到 1 之间生成它们，而不是从-1 到 1，因为在_PYTHAG宏中对它们进行平方运算将消除任何负值：

for(ULL i=0; i < iters; i++)
 {
     x = curand_uniform(&cr_state);
     y = curand_uniform(&cr_state);

     if(_PYTHAG(x,y) <= 1.0f)
         hits[tid]++;
 }

我们现在可以结束并关闭我们的内核，以及extern "C"块，最后用另一个}括号结束：

return;
}
}
''')

现在，让我们用get_function将 Python 包装函数传递给我们的内核。我们还将设置块和网格大小：每个块 32 个线程，每个网格 512 个块。让我们计算总线程数，并在 GPU 上设置一个数组来保存所有的命中（当然初始化为 0）：

pi_ker = ker.get_function("estimate_pi")
threads_per_block = 32
blocks_per_grid = 512
total_threads = threads_per_block * blocks_per_grid
hits_d = gpuarray.zeros((total_threads,),dtype=np.uint64)

让我们设置每个线程的迭代总数为 2²⁴：

iters = 2**24

我们现在可以像往常一样启动内核：

pi_ker(np.uint64(iters), hits_d, grid=(blocks_per_grid,1,1), block=(threads_per_block,1,1))

现在，让我们对数组中的命中次数求和，这给我们了总命中次数。让我们还计算数组中所有线程的总迭代次数：

total_hits = np.sum( hits_d.get() )
total = np.uint64(total_threads) * np.uint64(iters)

我们现在可以用Rational进行估计，就像这样：

est_pi_symbolic =  Rational(4)*Rational(int(total_hits), int(total) )

我们现在可以将其转换为浮点值：

est_pi = np.float(est_pi_symbolic.evalf())

让我们检查我们的估计与 NumPy 的常数值numpy.pi：

print "Our Monte Carlo estimate of Pi is : %s" % est_pi
print "NumPy's Pi constant is: %s " % np.pi
print "Our estimate passes NumPy's 'allclose' : %s" % np.allclose(est_pi, np.pi)

我们现在完成了。让我们从 IPython 中运行并检查一下（这个程序也可以在本书的存储库中的Chapter08下的monte_carlo_pi.py文件中找到）。

CUDA 数学 API

现在，我们将看一下CUDA 数学 API。这是一个库，由设备函数组成，类似于标准 C math.h库中的函数，可以从内核中的单个线程调用。这里的一个区别是，单精度和双精度浮点运算被重载，因此如果我们使用sin(x)，其中x是一个浮点数，sin 函数将产生一个 32 位浮点数作为输出，而如果x是一个 64 位双精度浮点数，那么sin的输出也将是一个 64 位值（通常，这是 32 位函数的正确名称，但它在末尾有一个f，比如sinf）。还有其他内在函数。内在函数是内置到 NVIDIA CUDA 硬件中的不太准确但更快的数学函数；通常，它们的名称与原始函数相似，只是在前面加上两个下划线—因此，内在的 32 位 sin 函数是__sinf。

明确积分的简要回顾

现在，我们将在 Python 中使用一些面向对象的编程，设置一个类，我们可以使用蒙特卡洛方法来评估函数的定积分。让我们停下来，谈谈我们的意思：假设我们有一个数学函数（就像你在微积分课上可能看到的那种类型），我们称之为f(x)。当我们在笛卡尔平面上在点a和b之间绘制它时，它可能看起来像这样：

现在，让我们仔细回顾一下定积分的确切含义——让我们将这个图中的第一个灰色区域表示为I，第二个灰色区域表示为II，第三个灰色区域表示为III。请注意，这里的第二个灰色区域是小于零的。这里的f的定积分，从a到b，将是值I - II + III，我们将在数学上表示为。一般来说，从a到b的定积分就是所有在f函数和 x 轴之间的总“正”区域的总和，其中 y > 0，减去所有在f函数和 x 轴之间的总“负”区域的总和，其中 y < 0，位于a和b之间。

有许多方法可以计算或估计两点之间函数的定积分。 在微积分课程中可能见过的一种方法是找到一个封闭形式的解：找到f的反导数F，并计算F(b) - F(a)。 然而，在许多领域，我们将无法找到确切的反导数，而必须以数值方式确定定积分。 这正是蒙特卡罗积分的想法：我们在a和b之间的许多随机点上评估f，然后使用这些点来估计定积分。

使用蒙特卡罗方法计算定积分

我们现在将使用 CUDA Math API 来表示任意数学函数f，同时使用 cuRAND 库来实现蒙特卡罗积分。 我们将使用元编程来实现这一点：我们将使用 Python 从代码模板生成设备函数的代码，这将插入适当的蒙特卡罗核心以进行积分。

这里的想法是它看起来和行为类似于我们在 PyCUDA 中看到的一些元编程工具，比如ElementwiseKernel。

让我们首先将适当的模块导入到我们的新项目中：

import pycuda.autoinit
import pycuda.driver as drv
from pycuda import gpuarray
from pycuda.compiler import SourceModule
import numpy as np

我们将使用 Python 中称为基于字典的字符串格式化的技巧。 让我们在继续之前花一分钟来了解一下。 假设我们正在编写一段 CUDA C 代码，并且我们不确定是否希望特定的变量集是 float 还是 double；也许看起来像这样：code_string="float x, y; float * z;"。 我们实际上可能希望格式化代码，以便我们可以随时在浮点数和双精度之间切换。 让我们将字符串中所有对float的引用更改为%(precision)s—code_string="%(precision)s x, y; %(precision)s * z;"。 现在，我们可以设置一个适当的字典，它将用double交换%(presision)s，即code_dict = {'precision' : 'double'}，并使用code_double = code_string % code_dict获取新的双精度字符串。 让我们看一下：

现在，让我们想一想我们想要我们的新蒙特卡洛积分器如何工作。 我们还将使其接受一个使用 CUDA Math API 编写的数学方程的字符串，以定义我们想要积分的函数。 然后，我们可以使用我们刚学到的字典技巧将这个字符串嵌入到代码中，并使用它来积分任意函数。 我们还将使用模板在float和double精度之间进行切换，根据用户的自由裁量。

我们现在可以开始我们的 CUDA C 代码：

MonteCarloKernelTemplate = '''
#include <curand_kernel.h>

我们将保留以前的无符号 64 位整数宏ULL。 让我们为 x 的倒数（_R）和平方（_P2）定义一些新的宏：

#define ULL unsigned long long
#define _R(z) ( 1.0f / (z) )
#define _P2(z) ( (z) * (z) )

现在，让我们定义一个设备函数，我们的方程字符串将插入其中。 当我们必须从字典中交换文本时，我们将使用math_function值。 我们将有另一个值称为p，表示精度（将是float或double）。 我们将称这个设备函数为f。 我们将在函数的声明中放置一个inline，这将节省我们一些时间，因为从核心调用时会有一些分支：

__device__ inline %(p)s f(%(p)s x)
{
    %(p)s y;
    %(math_function)s;
    return y;
}

现在，让我们考虑一下这将如何工作——我们声明一个名为y的 32 位或 64 位浮点值，调用math_function，然后返回y。 math_function，如果它是对输入参数x进行操作并将某个值设置为y的一些代码，那么这只有意义。 让我们记住这一点，然后继续。

我们现在将开始编写我们的蒙特卡洛积分核。 让我们记住，我们必须使用extern "C"关键字使我们的 CUDA 核可从普通 C 中访问。 然后我们将设置我们的核心。

首先，我们将用iters指示内核中每个线程应该采样多少随机样本；然后我们用lo指示积分的下界（b），用hi指示上界（a），并传入一个数组ys_out来存储每个线程的部分积分集合（我们稍后将对ys_out求和，以得到从主机端的lo到hi的完整定积分值）。再次注意我们将精度称为p：

extern "C" {
__global__ void monte_carlo(int iters, %(p)s lo, %(p)s hi, %(p)s * ys_out)
{

我们将需要一个curandState对象来生成随机值。我们还需要找到全局线程 ID 和线程的总数。由于我们正在处理一维数学函数，因此在一维x中设置我们的块和网格参数是有意义的：

curandState cr_state;
int tid = blockIdx.x * blockDim.x + threadIdx.x;
int num_threads = blockDim.x * gridDim.x;

现在我们将计算单个线程将处理的lo和hi之间的面积量。我们将通过将整个积分的长度（即hi - lo）除以线程的总数来实现这一点。:

再次注意我们如何使用模板技巧，以便这个值可以是多精度的。

%(p)s t_width = (hi - lo) / ( %(p)s ) num_threads;

回想一下，我们有一个名为iters的参数；这表示每个线程将采样多少个随机值。我们需要稍后知道样本的密度是多少；也就是说，每单位距离的平均样本数。我们可以这样计算，记得将整数iters强制转换为浮点值：

%(p)s density = ( ( %(p)s ) iters ) / t_width;

回想一下，我们正在将我们正在积分的区域按线程数进行划分。这意味着每个线程将有自己的起点和终点。由于我们正在为每个线程公平地划分长度，我们可以这样计算：

%(p)s t_lo = t_width*tid + lo;
 %(p)s t_hi = t_lo + t_width;

现在我们可以像之前一样初始化 cuRAND，确保每个线程都从自己的个体种子生成随机值：

curand_init( (ULL)  clock() + (ULL) tid, (ULL) 0, (ULL) 0, &cr_state);

在开始采样之前，我们需要设置一些额外的浮点值。y将保存从t_lo到t_hi的积分估计的最终值，y_sum将保存所有采样值的总和。我们还将使用rand_val变量来保存我们生成的原始随机值，x来存储我们将要从中采样的区域的缩放随机值：

%(p)s y, y_sum = 0.0f;
%(p)s rand_val, x;

现在，让我们循环从我们的函数中取样值，将这些值加入y_sum中。需要注意的一点是curand_uniform末尾的%(p_curand)s——这个函数的 32 位浮点版本是curand_uniform，而 64 位版本是curand_uniform_double。稍后我们将根据这里使用的精度级别，用_double或空字符串来替换它。还要注意我们如何缩放rand_val，使得x落在t_lo和t_hi之间，记住 cuRAND 中的随机均匀分布只产生 0 到 1 之间的值：

for (int i=0; i < iters; i++)
{
    rand_val = curand_uniform%(p_curand)s(&cr_state);
    x = t_lo + t_width * rand_val;
    y_sum += f(x);
}

现在我们可以通过密度将y_sum除以t_hi到t_lo的子积分的值：

y = y_sum / density;

我们将这个值输出到数组中，并关闭我们的 CUDA 内核，以及extern "C"，以及最终的闭括号。我们已经写完了 CUDA C，所以我们将用三个引号结束这一部分：

ys_out[tid] = y;
}
}
'''

现在我们要做一些不同的事情——我们将设置一个类来处理我们的定积分。让我们称之为MonteCarloIntegrator。当然，我们将首先编写构造函数，也就是__init__函数。这是我们将输入对象引用self的地方。让我们将math_function的默认值设置为'y = sin(x)'，默认精度为'd'，即双精度。我们还将把lo的默认值设置为 0，hi的默认值设置为π的 NumPy 近似值。最后，我们将有每个线程将采样的随机样本数（samples_per_thread）和我们将在其上启动内核的网格大小（num_blocks）的值。

让我们从将文本字符串math_function存储在self对象中开始这个函数，以便以后使用：

def __init__(self, math_function='y = sin(x)', precision='d', lo=0, hi=np.pi, samples_per_thread=10**5, num_blocks=100):

        self.math_function = math_function

现在，让我们设置与我们选择的浮点精度相关的值，这将在以后设置我们的模板字典时特别需要，特别是为了存储lo和hi的值。让我们还在对象中存储lo和hi的值。如果用户输入无效的数据类型，或者hi实际上小于lo，让我们确保引发异常错误：

         if precision in [None, 's', 'S', 'single', np.float32]:
             self.precision = 'float'
             self.numpy_precision = np.float32
             self.p_curand = ''
         elif precision in ['d','D', 'double', np.float64]:
             self.precision = 'double'
             self.numpy_precision = np.float64
             self.p_curand = '_double'
         else:
             raise Exception('precision is invalid datatype!')

     if (hi - lo <= 0):
         raise Exception('hi - lo <= 0!')
     else:
         self.hi = hi
         self.lo = lo

现在，我们可以设置我们的代码模板字典：

MonteCarloDict = {'p' : self.precision, 'p_curand' : self.p_curand, 'math_function' : self.math_function}

现在我们可以使用基于字典的字符串格式生成实际的最终代码，并进行编译。让我们还通过在SourceModule中设置options=['-w']来关闭nvcc编译器的警告：

self.MonteCarloCode = MonteCarloKernelTemplate % MonteCarloDict

self.ker = SourceModule(no_extern_c=True , options=['-w'], source=self.MonteCarloCode)

现在，我们将在我们的对象中设置一个函数引用到我们编译的内核，使用get_function。在继续之前，让我们保存对象中的剩余两个参数：

self.f = self.ker.get_function('monte_carlo')
self.num_blocks = num_blocks
self.samples_per_thread = samples_per_thread

现在，虽然我们需要不同实例的MonteCarloIntegrator对象来评估不同数学函数或浮点精度的定积分，但我们可能希望在不同的lo和hi边界上评估相同的积分，改变线程/网格大小，或者改变我们在每个线程上取样的数量。幸运的是，这些都是容易进行的修改，并且都可以在运行时进行。

我们将设置一个特定的函数来评估给定对象的积分。我们将这些参数的默认值设置为在调用构造函数时存储的值：

def definite_integral(self, lo=None, hi=None, samples_per_thread=None, num_blocks=None):
    if lo is None or hi is None:
        lo = self.lo
        hi = self.hi
    if samples_per_thread is None:
        samples_per_thread = self.samples_per_thread
    if num_blocks is None:
        num_blocks = self.num_blocks
        grid = (num_blocks,1,1)
    else:
        grid = (num_blocks,1,1)

    block = (32,1,1)
    num_threads = 32*num_blocks

我们可以通过设置一个空数组来存储部分子积分并启动内核来完成这个函数。然后我们需要对子积分求和以获得最终值，然后返回：

self.ys = gpuarray.empty((num_threads,) , dtype=self.numpy_precision)

self.f(np.int32(samples_per_thread), self.numpy_precision(lo), self.numpy_precision(hi), self.ys, block=block, grid=grid)

self.nintegral = np.sum(self.ys.get() )

return np.sum(self.nintegral)

我们已经准备好尝试这个了。让我们只是设置一个具有默认值的类——这将从 0 到π积分y = sin(x)。如果您记得微积分，sin(x)的反导数是-cos(x)，所以我们可以这样评估定积分：

因此，我们应该得到一个接近 2 的数值。让我们看看我们得到了什么：

编写一些测试用例

现在，我们终于将看到如何使用 CUDA Math API 通过math_function参数编写一些测试用例来测试我们的类。如果您有 C/C++标准数学库的经验，这将会相当简单。同样，这些函数是重载的，这样当我们在单精度和双精度之间切换时，我们就不必更改任何名称。

我们已经看到了一个例子，即y = sin(x)。让我们尝试一些更有雄心的东西：

我们将从a=11.733 积分到b=18.472，然后检查我们的蒙特卡洛积分器的输出与另一个来源的已知值进行比较。在这里，Mathematica 指出这个定积分的值是 8.9999，所以我们将与其进行比较。

现在，让我们考虑如何表示这个函数：这里，log指的是自然对数（也称为ln），在 Math API 中就是log(x)。我们已经设置了一个宏来表示平方，所以我们可以用_P2(sin(x))来表示sin²(x)。现在我们可以用y = log(x)*_P2(sin(x))来表示整个函数。

让我们使用以下方程，从a=.9积分到b=4：

记住，_R是我们设置的倒数宏，我们可以这样用 Math API 来写这个函数：

'y = _R( 1 + sinh(2*x)*_P2(log(x)) )' 

在我们继续之前，让我们注意到 Mathematica 告诉我们这个定积分的值是.584977。

让我们再检查一个函数。让我们有点雄心勃勃地说它是这样的：

我们可以将这表示为'y = (cosh(x)*sin(x))/ sqrt( pow(x,3) + _P2(sin(x)))'；自然地，sqrt是分母中的平方根，pow允许我们取任意幂的值。当然，sin(x)是sin(x)，cosh(x)是cosh(x)。我们从a=1.85 积分到b=4.81；Mathematica 告诉我们这个积分的真实值是-3.34553。

我们现在准备检查一些测试用例，并验证我们的蒙特卡洛积分是否有效！让我们迭代一个列表，其第一个值是指示函数（使用 Math API）的字符串，第二个值指示积分的下限，第三个值指示积分的上限，最后一个值指示用 Mathematica 计算出的预期值：

if __name__ == '__main__':

    integral_tests = [('y =log(x)*_P2(sin(x))', 11.733 , 18.472, 8.9999), ('y = _R( 1 + sinh(2*x)*_P2(log(x)) )', .9, 4, .584977), ('y = (cosh(x)*sin(x))/ sqrt( pow(x,3) + _P2(sin(x)))', 1.85, 4.81, -3.34553) ]

我们现在可以迭代这个列表，看看我们的算法与 Mathematica 相比效果如何：

for f, lo, hi, expected in integral_tests:
    mci = MonteCarloIntegrator(math_function=f, precision='d', lo=lo, hi=hi)
    print 'The Monte Carlo numerical integration of the function\n \t f: x -> %s \n \t from x = %s to x = %s is : %s ' % (f, lo, hi, mci.definite_integral())
    print 'where the expected value is : %s\n' % expected

现在就运行这个：

这也可以在本书存储库的Chapter08目录下的monte_carlo_integrator.py文件中找到。

CUDA Thrust 库

我们现在将看一下 CUDA Thrust 库。这个库的核心特性是一个高级向量容器，类似于 C++自己的向量容器。虽然这听起来可能很琐碎，但这将使我们在 CUDA C 中编程时更少地依赖指针、malloc 和 free。与 C++向量容器一样，Thrust 的向量容器会自动处理元素的调整大小和连接，并且借助 C++析构函数的魔力，释放也会在 Thrust 向量对象超出范围时自动处理。

Thrust 实际上提供了两个向量容器：一个用于主机端，一个用于设备端。主机端的 Thrust 向量与 STL 向量几乎相同，主要区别在于它可以更轻松地与 GPU 交互。让我们用适当的 CUDA C 代码写一点代码，以了解它是如何工作的。

让我们从包含语句开始。我们将使用主机和设备端向量的头文件，并且还将包括 C++的iostream库，这将允许我们在终端上执行基本的 I/O 操作：

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <iostream>

让我们只使用标准的 C++命名空间（这样我们在检查输出时就不必输入std::分辨率运算符）：

using namespace std;

我们现在将制作我们的主函数，并在主机端设置一个空的 Thrust 向量。同样，这些都是 C++模板，所以我们必须在声明时使用< >括号选择数据类型。我们将把这设置为一个整数数组：

int main(void)
{
 thrust::host_vector<int> v;

现在，让我们通过使用push_back在v的末尾添加一些整数，就像我们用常规 STL 向量一样：

v.push_back(1);
v.push_back(2);
v.push_back(3);
v.push_back(4);

我们现在将迭代向量中的所有值，并输出每个值：

这里的输出应该是v[0] == 1到v[3] == 4。

for (int i = 0; i < v.size(); i++)
    cout << "v[" << i << "] == " << v[i] << endl;

到目前为止，这可能看起来很琐碎。让我们在 GPU 上设置一个 Thrust 向量，然后将内容从v复制过去：

thrust::device_vector<int> v_gpu = v;

是的，就这样了——只有一行，我们就完成了。现在主机上的v的所有内容都将被复制到设备上的v_gpu！（如果这让你感到惊讶，请再看一下第六章，调试和分析您的 CUDA 代码，想想在这之前我们需要多少行。）

让我们尝试在我们的新 GPU 向量上使用push_back，看看我们是否可以将另一个值连接到它上面：

v_gpu.push_back(5);

我们现在将检查v_gpu的内容，如下所示：

for (int i = 0; i < v_gpu.size(); i++)
    std::cout << "v_gpu[" << i << "] == " << v_gpu[i] << std::endl;

这部分应该输出v_gpu[0] == 1到v_gpu[4] == 5。

再次感谢这些对象的析构函数，我们不必进行任何清理工作，释放任何已分配内存的块。现在我们可以从程序中返回，我们完成了：

    return 0;
}

在 Thrust 中使用函数对象

让我们看看如何在 Thrust 中使用称为functors的概念。在 C++中，functor是一个看起来和行为像函数的类或结构对象；这让我们可以使用看起来和行为像函数的东西，但可以保存一些不必每次使用时都设置的参数。

让我们用适当的包含语句开始一个新的 Thrust 程序，并使用标准命名空间：

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <iostream>
using namespace std;

现在，让我们设置一个基本的函数对象。我们将使用struct来表示这个函数对象，而不是class。这将是一个加权乘法函数，我们将在一个名为w的浮点数中存储权重。我们将创建一个构造函数，用默认值1设置权重：

struct multiply_functor {
 float w;
 multiply_functor(float _w = 1) : w(_w) {}

现在，我们将使用operator()关键字设置我们的函数对象；这将告诉编译器将以下代码块视为此类型对象的default函数。请记住，这将在 GPU 上作为设备函数运行，因此我们在整个代码块之前加上__device__。我们用括号指示输入，并输出适当的值，这只是一个缩放的倍数。现在，我们可以用};关闭我们的结构的定义：

    __device__ float operator() (const float & x, const float & y) { 
        return w * x * y;
     }
};

现在，让我们使用这个来制作一个基本的点积函数；回想一下，这需要两个数组之间的逐点乘法，然后是一个reduce类型的求和。让我们首先声明我们的函数，并创建一个新的向量z，它将保存逐点乘法的值：

float dot_product(thrust::device_vector<float> &v, thrust::device_vector<float> &w ), thrust::device_vector<float> &z)
{
 thrust::device_vector<float> z(v.size());

我们现在将使用 Thrust 的transform操作，它将对v和w的输入进行逐点操作，并输出到z。请注意，我们将函数对象输入到 transform 的最后一个槽中；通过这样使用普通的闭括号，它将使用构造函数的默认值（w = 1），因此这将作为一个普通的、非加权/缩放的点积：

thrust::transform(v.begin(), v.end(), w.begin(), z.begin(), multiply_functor());

我们现在可以使用 Thrust 的 reduce 函数对z进行求和。让我们返回这个值：

return thrust::reduce(z.begin(), z.end());
}

我们完成了。现在，让我们编写一些测试代码——我们将只计算向量[1,2,3]和[1,1,1]的点积，这对我们来说很容易检查。（结果将是 6。）

让我们使用push_back设置第一个向量v：

int main(void)
{
    thrust::device_vector<float> v;
    v.push_back(1.0f);
    v.push_back(2.0f);
    v.push_back(3.0f);

现在，我们可以声明一个大小为3的向量w，并使用 Thrust 的 fill 函数将其默认值设置为1，如下所示：

thrust::device_vector<float> w(3);
thrust::fill(w.begin(), w.end(), 1.0f);

让我们进行一次检查，确保我们的值被正确设置，通过将它们的值输出到cout：

for (int i = 0; i < v.size(); i++)
 cout << "v[" << i << "] == " << v[i] << endl;

for (int i = 0; i < w.size(); i++)
 cout << "w[" << i << "] == " << w[i] << endl;

现在，我们可以检查我们的点积的输出，然后从程序中返回：

cout << "dot_product(v , w) == " << dot_product(v,w) << endl;
return 0;
}

让我们编译这个程序（在 Linux 或 Windows 的命令行中使用nvcc thrust_dot_product.cu -o thrust_dot_product）并运行它：

这个代码也可以在本书存储库的Chapter08目录中的thrust_dot_product.cu文件中找到。

摘要

在本章中，我们看了如何通过选择适当的种子在 cuRAND 中初始化随机数流。由于计算机是确定性设备，它们只能生成伪随机数列表，因此我们的种子应该是真正随机的；通常，将线程 ID 添加到毫秒级的时钟时间中就足够满足大多数目的。

然后，我们看了如何使用 cuRAND 的均匀分布来对 Pi 进行基本估计。然后，我们承担了一个更有雄心的项目，创建一个可以计算任意函数的定积分的 Python 类；我们使用了一些元编程的思想，结合 CUDA Math API 来定义这些任意函数。最后，我们简要介绍了 CUDA Thrust 库，它通常用于在 Python 之外编写纯 CUDA C 程序。Thrust 最显著的提供了一个类似于标准 C++ vector的device_vector容器。这减少了在 CUDA C 中使用指针的一些认知开销。

最后，我们简要介绍了如何使用 Thrust 和适当的函数对象来执行简单的point-wise和reduce操作，即实现一个简单的点积函数。

问题

1. 尝试重写蒙特卡洛积分示例（在monte_carlo_integrator.py的__main__函数中）以使用 CUDA 的intrinsic函数。精度与以前相比如何？

2. 在我们所有的 cuRAND 示例中，我们只使用了均匀分布。你能说出在 GPU 编程中使用正态（高斯）随机分布的一个可能的用途或应用吗？

3. 假设我们使用两个不同的种子生成一个包含 100 个伪随机数的列表。我们应该将它们连接成一个包含 200 个数字的列表吗？

4. 在最后一个示例中，在multiply_functor结构体的operator()函数定义之前添加__host__，现在，看看是否可以直接使用这个函数对象实现一个主机端的点积函数，而不需要任何进一步的修改。

5. 看一下 Thrust examples目录中的strided_range.cu文件。你能想到如何使用这个来使用 Thrust 实现通用矩阵乘法吗？

6. 在定义一个函数对象时，operator()函数的重要性是什么？

第九章：实现深度神经网络

我们现在将利用我们对 GPU 编程的积累知识，使用 PyCUDA 实现我们自己的深度神经网络（DNN）。在过去的十年里，DNN 吸引了很多关注，因为它们为机器学习（ML）提供了一个强大而优雅的模型。DNN 也是第一个能够展示 GPU 真正强大之处的应用之一（除了渲染图形），通过利用它们的大规模并行吞吐量，最终帮助 NVIDIA 成为人工智能领域的主要参与者。

在本书的过程中，我们大多是按章节覆盖单个主题的——在这里，我们将在我们学到的许多主题的基础上构建我们自己的 DNN 实现。虽然目前有几个面向普通公众的基于 GPU 的 DNN 的开源框架，例如 Google 的 TensorFlow 和 Keras，微软的 CNTK，Facebook 的 Caffe2 和 PyTorch，但从头开始实现一个 DNN 非常有教育意义，这将使我们更深入地了解和欣赏 DNN 所需的基础技术。我们有很多材料要涵盖，所以在简要介绍一些基本概念后，我们将直奔主题。

在本章中，我们将研究以下内容：

• 理解人工神经元（AN）是什么

• 理解如何将多个 AN 组合在一起形成深度神经网络（DNN）

• 在 CUDA 和 Python 中从头开始实现 DNN

• 理解如何使用交叉熵损失来评估神经网络的输出

• 实现梯度下降来训练 NN

• 学习如何在小数据集上训练和测试 NN

技术要求

本章需要一台装有现代 NVIDIA GPU（2016 年以后）的 Linux 或 Windows 10 PC，并安装了所有必要的 GPU 驱动程序和 CUDA Toolkit（9.0 及以上）。还需要一个合适的 Python 2.7 安装（如 Anaconda Python 2.7），并安装了 PyCUDA 模块。

本章的代码也可以在 GitHub 上找到：github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA。

有关本章的先决条件的更多信息，请查看本书的前言。有关软件和硬件要求，请查看github.com/PacktPublishing/Hands-On-GPU-Programming-with-Python-and-CUDA中的 README 文件。

人工神经元和神经网络

让我们简要回顾一些机器学习（ML）和神经网络（NNs）的基础知识。在机器学习中，我们的目标是使用具有特定标记类别或特征的数据集，并利用这些示例来训练我们的系统以预测未来数据的值。我们称根据先前训练数据预测未来数据的类别或标签的程序或函数为分类器。

有许多类型的分类器，但在这里我们将专注于 NNs。NNs 的理念是它们（据说）以类似于人脑的方式工作，通过使用一组人工神经元（ANs）来学习和分类数据，所有这些神经元连接在一起形成特定的结构。不过，让我们暂停一下，看看一个单独的 AN 是什么。在数学上，这只是从线性空间R^n到R的仿射函数，如下所示：

我们可以看到这可以被描述为一个常量权重向量w和输入向量x之间的点积，最后加上一个额外的偏置常量b。（再次强调，这个函数的唯一输入是x；其他值都是常数！）

现在，单个 AN 本身是相当无用（而且愚蠢）的，只有当它们与大量其他 AN 合作时，它们的智能才会显现出来。我们的第一步是将一系列相似的 AN 堆叠在一起，以形成我们将称之为密集层（DL）的东西。这是密集的，因为每个神经元将处理来自x的每个输入值 - 每个 AN 将接收来自R^{n**的数组或向量值，并在**R**中输出一个值。由于有*m*个神经元，这意味着它们的输出集体位于**R}m空间中。我们将注意到，如果我们堆叠我们层中每个神经元的权重，以形成一个m x n的权重矩阵，然后我们可以通过矩阵乘法计算每个神经元的输出，然后加上适当的偏差：

现在，假设我们想要构建一个能够对k个不同类别进行分类的 NN 分类器；我们可以创建一个新的附加密集层，该层接收来自先前密集层的m个值，并输出k个值。假设我们对每一层都有适当的权重和偏差值（这显然不容易找到），并且在每一层之后也有适当的激活函数设置（我们稍后会定义），这将作为我们k个不同类别之间的分类器，根据最终层的输出给出x落入每个类别的概率。当然，我们在这里走得太远了，但这就是 NN 的工作原理。

现在，似乎我们可以将密集层连接到长链中以实现分类。这就是所谓的 DNN。当我们有一层不直接连接到输入或输出时，这就是一个隐藏层。DNN 的优势在于额外的层允许 NN 捕捉浅层 NN 无法捕捉到的数据的抽象和细微差别。

实现人工神经元的密集层

现在，让我们实现 NN 最重要的构建模块，密集层。让我们从声明 CUDA 内核开始，就像这样：

__global__ void dense_eval(int num_outputs, int num_inputs, int relu, int sigmoid, float * w, float * b, float * x, float *y, int batch_size, int w_t, int b_t, float delta)

让我们逐个检查输入。num_outputs当然表示这一层的总输出数量；这正好是该层的神经元数量。num_inputs告诉我们输入数据的大小。为relu和sigmoid设置正值将表示我们应该在该层的输出上使用相应的激活函数，我们稍后会定义。w和b是包含该层权重和偏差的数组，而x和y将作为我们的输入和输出。通常，我们希望一次对多个数据进行分类。我们可以通过将batch_size设置为我们希望预测的点数来指示这一点。最后，w_t、b_t和delta将在训练过程中使用，通过梯度下降来确定该层的适当权重和偏差。（我们将在后面的部分中更多地了解梯度下降。）

现在，让我们开始编写我们的内核。我们将在每个输出上并行计算，因此我们将设置一个整数i作为全局线程 ID，然后使用适当的if语句让任何不必要的额外线程不执行任何操作：

{
 int i = blockDim.x*blockIdx.x + threadIdx.x;

 if (i < num_outputs)
 {

现在，让我们用适当的for循环迭代批处理中的每个数据点：

for(int k=0; k < batch_size; k++)
 { 

我们将从权重和输入中的 32 位浮点数中进行乘法和累加，得到 64 位双精度temp，然后加上适当的偏差点。然后我们将这个值强制转换回 32 位浮点数并将值放入输出数组中，然后关闭对k的循环：

double temp = 0.0f;
 for (int j = 0; j < num_inputs; j++)
 {
   temp += ((double) w[(num_inputs)*i + j ] ) * ( (double) x[k*num_inputs + j]);
 }
 temp += (double) b[i];
 y[k * num_outputs + i] = (float) temp;  
}

乘法和累加类型的操作通常会导致严重的数值精度损失。这可以通过使用更高精度的临时变量来存储操作过程中的值，然后在操作完成后将此变量强制转换回原始精度来减轻。

要训练一个 NN，我们最终将不得不计算我们的 NN 对每个权重和偏差的导数（来自微积分），这是针对特定输入批次的。请记住，数学函数f在值x处的导数可以估计为f**(x + δ) - f(x) / δ，其中 delta（δ）是某个足够小的正值。我们将使用输入值w_t和b_t来指示内核是否要计算相对于特定权重或偏差的导数；否则，我们将将这些输入值设置为负值，仅对此层进行评估。我们还将设置 delta 为适当小的值，用于计算导数，并使用它来增加适当偏差或权重的值：

if( w_t >= 0 && i == (w_t / num_inputs))
 {
 int j = w_t % num_inputs;
 for(int k=0; k < batch_size; k++)
  y[k*num_outputs + i] += delta*x[k*num_inputs+j];
}
if( b_t >= 0 && i == b_t )
 {
  for(int k=0; k < batch_size; k++)
  y[k*num_outputs + i] += delta;
 }

现在，我们将添加一些代码，用于所谓的修正线性单元（或ReLU）和Sigmoid 激活函数。这些用于处理密集神经层的直接输出。ReLU 只将所有负值设为 0，同时作为正输入的恒等式，而 sigmoid 只计算每个值上的sigmoid函数的值（1 / (1 + e^(-x))）。ReLU（或任何其他激活函数）用于 NN 中隐藏层之间，作为使整个 NN 成为非线性函数的手段；否则，整个 NN 将构成一个微不足道（且计算效率低下）的矩阵操作。（虽然可以在层之间使用许多其他非线性激活函数，但发现 ReLU 对训练来说是一个特别有效的函数。）Sigmoid 用作 NN 中用于标签的最终层，即可能为给定输入分配多个标签的层，而不是将输入分配给单个类别。

让我们在文件中稍微上移一点，甚至在我们开始定义这个 CUDA 内核之前，将这些操作定义为 C 宏。我们还要记得在我们写完 CUDA-C 代码后将其放入其中：

DenseEvalCode = '''
#define _RELU(x) ( ((x) > 0.0f) ? (x) : 0.0f )
#define _SIGMOID(x) ( 1.0f / (1.0f + expf(-(x)) ))

现在，我们将使用内核输入relu和sigmoid来指示我们是否应该使用这些额外的层；我们将从这些中获取积极的输入，以指示它们应该分别使用。我们可以添加这个，关闭我们的内核，并将其编译成可用的 Python 函数：

if(relu > 0 || sigmoid > 0)
for(int k=0; k < batch_size; k++)
 { 
   float temp = y[k * num_outputs + i];
   if (relu > 0)
    temp = _RELU(temp);
   if (sigmoid > 0)
    temp = _SIGMOID(temp);
   y[k * num_outputs + i] = temp; 
  }
 }
 return;
}
'''
eval_mod = SourceModule(DenseEvalCode)
eval_ker = eval_mod.get_function('dense_eval')

现在，让我们回到文件的开头，并设置适当的导入语句。请注意，我们将包括csv模块，该模块将用于处理测试和训练的数据输入：

from __future__ import division
import pycuda.autoinit
import pycuda.driver as drv
from pycuda import gpuarray
from pycuda.compiler import SourceModule
from pycuda.elementwise import ElementwiseKernel
import numpy as np
from Queue import Queue
import csv
import time

现在，让我们继续设置我们的密集层；我们将希望将其包装在一个 Python 类中以便使用，这将使我们在开始将这些密集层连接成一个完整的 NN 时更加轻松。我们将称之为class DenseLayer并开始编写构造函数。这里的大部分输入和设置应该是不言自明的：我们绝对应该添加一个选项，从预训练的网络中加载权重和偏差，并且我们还将包括指定默认delta值以及默认流的选项。（如果没有给出权重或偏差，权重将被初始化为随机值，而所有偏差将设置为 0。）我们还将在这里指定是否使用 ReLU 或 sigmoid 层。最后，请注意我们如何设置块和网格大小：

class DenseLayer:
    def __init__(self, num_inputs=None, num_outputs=None, weights=None, b=None, stream=None, relu=False, sigmoid=False, delta=None):
        self.stream = stream

        if delta is None:
            self.delta = np.float32(0.001)
        else:
            self.delta = np.float32(delta)

        if weights is None:
            weights = np.random.rand(num_outputs, num_inputs) - .5
            self.num_inputs = np.int32(num_inputs)
        self.num_outputs = np.int32(num_outputs) 

        if type(weights) != pycuda.gpuarray.GPUArray:
            self.weights = gpuarray.to_gpu_async(np.array(weights, 
            dtype=np.float32) , stream = self.stream)
        else:
            self.weights = weights

        if num_inputs is None or num_outputs is None:
            self.num_inputs = np.int32(self.weights.shape[1])
            self.num_outputs = np.int32(self.weights.shape[0])

        else:
            self.num_inputs = np.int32(num_inputs)
            self.num_outputs = np.int32(num_outputs)

        if b is None:
            b = gpuarray.zeros((self.num_outputs,),dtype=np.float32)

        if type(b) != pycuda.gpuarray.GPUArray:
            self.b = gpuarray.to_gpu_async(np.array(b, 
            dtype=np.float32) , stream = self.stream)
        else:
            self.b = b 

        self.relu = np.int32(relu)
        self.sigmoid = np.int32(sigmoid)

        self.block = (32,1,1)
        self.grid = (int(np.ceil(self.num_outputs / 32)), 1,1)

现在，我们将在这个类中设置一个函数来评估来自这一层的输入；我们将仔细检查输入（x）以确定它是否已在 GPU 上（如果没有，则将其转移到gpuarray），并让用户指定预分配的gpuarray用于输出（y），如果没有指定，则手动分配一个输出数组。我们还将检查训练时的 delta 和w_t/b_t值，以及batch_size。然后我们将在x输入上运行核函数，输出到y，最后将y作为输出值返回：

def eval_(self, x, y=None, batch_size=None, stream=None, delta=None, w_t = None, b_t = None):

if stream is None:
    stream = self.stream

if type(x) != pycuda.gpuarray.GPUArray:
    x = gpuarray.to_gpu_async(np.array(x,dtype=np.float32), stream=self.stream)

if batch_size is None:
    if len(x.shape) == 2:
        batch_size = np.int32(x.shape[0])
    else:
        batch_size = np.int32(1)

if delta is None:
    delta = self.delta

delta = np.float32(delta)

if w_t is None:
    w_t = np.int32(-1)

if b_t is None:
    b_t = np.int32(-1)

if y is None:
    if batch_size == 1:
        y = gpuarray.empty((self.num_outputs,), dtype=np.float32)
    else:
        y = gpuarray.empty((batch_size, self.num_outputs), dtype=np.float32)

    eval_ker(self.num_outputs, self.num_inputs, self.relu, self.sigmoid, self.weights, self.b, x, y, np.int32(batch_size), w_t, b_t, delta , block=self.block, grid=self.grid , stream=stream)

 return y

现在，我们已经完全实现了一个密集层！

实现 softmax 层

现在，让我们看看如何实现softmax 层。正如我们已经讨论过的，sigmoid 层用于为类分配标签——也就是说，如果您想要从输入中推断出多个非排他性特征，您应该使用 sigmoid 层。softmax 层用于仅通过推断为样本分配单个类别——这是通过计算每个可能类别的概率来实现的（当然，所有类别的概率总和为 100%）。然后我们可以选择具有最高概率的类别来给出最终分类。

现在，让我们看看 softmax 层到底做了什么——给定一组N个实数(c[0], ..., c[N-1])，我们首先计算每个数字的指数函数的总和，然后计算每个数字除以这个总和的指数，得到 softmax：

让我们从我们的实现开始。我们将首先编写两个非常简短的 CUDA 内核：一个用于计算每个输入的指数，另一个用于计算所有点的平均值：

SoftmaxExpCode='''
__global__ void softmax_exp( int num, float *x, float *y, int batch_size)
{
 int i = blockIdx.x * blockDim.x + threadIdx.x;

 if (i < num)
 {
  for (int k=0; k < batch_size; k++)
  {
   y[num*k + i] = expf(x[num*k+i]);
  }
 }
}
'''
exp_mod = SourceModule(SoftmaxExpCode)
exp_ker = exp_mod.get_function('softmax_exp')

SoftmaxMeanCode='''
__global__ void softmax_mean( int num, float *x, float *y, int batch_size)
{
 int i = blockDim.x*blockIdx.x + threadIdx.x;

 if (i < batch_size)
 {
  float temp = 0.0f;

  for(int k=0; k < num; k++)
   temp += x[i*num + k];

  for(int k=0; k < num; k++)
   y[i*num+k] = x[i*num+k] / temp;
 }

 return;
}'''

mean_mod = SourceModule(SoftmaxMeanCode)
mean_ker = mean_mod.get_function('softmax_mean')

现在，让我们编写一个 Python 包装类，就像我们之前做的那样。首先，我们将从构造函数开始，并使用num指示输入和输出的数量。如果需要，我们还可以指定默认流：

class SoftmaxLayer:
    def __init__(self, num=None, stream=None):
     self.num = np.int32(num)
     self.stream = stream

现在，让我们编写一个类似于密集层的eval_函数：

def eval_(self, x, y=None, batch_size=None, stream=None):
 if stream is None:
 stream = self.stream

 if type(x) != pycuda.gpuarray.GPUArray:
  temp = np.array(x,dtype=np.float32)
  x = gpuarray.to_gpu_async( temp , stream=stream)

 if batch_size==None:
  if len(x.shape) == 2:
   batch_size = np.int32(x.shape[0])
  else:
   batch_size = np.int32(1)
 else:
  batch_size = np.int32(batch_size)

 if y is None:
  if batch_size == 1:
   y = gpuarray.empty((self.num,), dtype=np.float32)
 else:
  y = gpuarray.empty((batch_size, self.num), dtype=np.float32)

 exp_ker(self.num, x, y, batch_size, block=(32,1,1), grid=(int( np.ceil( self.num / 32) ), 1, 1), stream=stream)

 mean_ker(self.num, y, y, batch_size, block=(32,1,1), grid=(int( np.ceil( batch_size / 32)), 1,1), stream=stream)

 return y

实现交叉熵损失

现在，让我们实现所谓的交叉熵损失函数。这用于在训练过程中测量神经网络在数据子集上的准确性；损失函数输出的值越大，我们的神经网络在正确分类给定数据方面就越不准确。我们通过计算期望输出和实际输出之间的标准平均对数熵差来实现这一点。为了数值稳定性，我们将限制输出值为1：

MAX_ENTROPY = 1

def cross_entropy(predictions=None, ground_truth=None):

 if predictions is None or ground_truth is None:
  raise Exception("Error! Both predictions and ground truth must be float32 arrays")

 p = np.array(predictions).copy()
 y = np.array(ground_truth).copy()

 if p.shape != y.shape:
  raise Exception("Error! Both predictions and ground_truth must have same shape.")

 if len(p.shape) != 2:
  raise Exception("Error! Both predictions and ground_truth must be 2D arrays.")

 total_entropy = 0

 for i in range(p.shape[0]):
  for j in range(p.shape[1]):
   if y[i,j] == 1: 
    total_entropy += min( np.abs( np.nan_to_num( np.log( p[i,j] ) ) ) , MAX_ENTROPY) 
   else: 
    total_entropy += min( np.abs( np.nan_to_num( np.log( 1 - p[i,j] ) ) ), MAX_ENTROPY)

 return total_entropy / p.size

实现一个顺序网络

现在，让我们实现最后一个类，将多个密集层和 softmax 层对象组合成一个连贯的前馈顺序神经网络。这将作为另一个类来实现，它将包含其他类。让我们首先从编写构造函数开始——我们将能够在这里设置最大批处理大小，这将影响为此网络分配多少内存——我们将在列表变量network_mem中存储一些用于权重和每个层的输入/输出的分配内存。我们还将在列表 network 中存储 DenseLayer 和 SoftmaxLayer 对象，并在 network_summary 中存储有关 NN 中每个层的信息。请注意，我们还可以在这里设置一些训练参数，包括 delta，用于梯度下降的流的数量（稍后我们将看到），以及训练时期的数量。

我们还可以看到开始时的另一个输入称为 layers。在这里，我们可以通过描述每个层来指示 NN 的构造，构造函数将通过迭代 layers 的每个元素并调用add_layer方法来创建它，接下来我们将实现这个方法：

class SequentialNetwork:
 def __init__(self, layers=None, delta=None, stream = None, max_batch_size=32, max_streams=10, epochs = 10):

 self.network = []
 self.network_summary = []
 self.network_mem = []

 if stream is not None:
  self.stream = stream
 else:
  self.stream = drv.Stream()

 if delta is None:
  delta = 0.0001

 self.delta = delta
 self.max_batch_size=max_batch_size
 self.max_streams = max_streams
 self.epochs = epochs

 if layers is not None:
  for layer in layers:
   add_layer(self, layer)

现在，让我们实现add_layer方法。我们将使用字典数据类型将有关该层的所有相关信息传递给顺序网络，包括层的类型（密集、softmax 等）、输入/输出的数量、权重和偏差。这将向对象的网络和network_summary列表变量追加适当的对象和信息，并适当地为network_mem列表分配gpuarray对象：

def add_layer(self, layer):
 if layer['type'] == 'dense':
  if len(self.network) == 0:
   num_inputs = layer['num_inputs']
  else:
   num_inputs = self.network_summary[-1][2]

  num_outputs = layer['num_outputs']
  sigmoid = layer['sigmoid']
  relu = layer['relu']
  weights = layer['weights']
  b = layer['bias']

  self.network.append(DenseLayer(num_inputs=num_inputs, num_outputs=num_outputs, sigmoid=sigmoid, relu=relu, weights=weights, b=b))
  self.network_summary.append( ('dense', num_inputs, num_outputs))

  if self.max_batch_size > 1:
   if len(self.network_mem) == 0:
self.network_mem.append(gpuarray.empty((self.max_batch_size, self.network_summary[-1][1]), dtype=np.float32))
 self.network_mem.append(gpuarray.empty((self.max_batch_size, self.network_summary[-1][2] ), dtype=np.float32 ) ) 
 else:
 if len(self.network_mem) == 0:
 self.network_mem.append( gpuarray.empty( (self.network_summary[-1][1], ), dtype=np.float32 ) )
 self.network_mem.append( gpuarray.empty((self.network_summary[-1][2], ), dtype=np.float32 ) ) 

 elif layer['type'] == 'softmax':

  if len(self.network) == 0:
   raise Exception("Error! Softmax layer can't be first!")

  if self.network_summary[-1][0] != 'dense':
   raise Exception("Error! Need a dense layer before a softmax layer!")

  num = self.network_summary[-1][2]
  self.network.append(SoftmaxLayer(num=num))
  self.network_summary.append(('softmax', num, num))

  if self.max_batch_size > 1:
   self.network_mem.append(gpuarray.empty((self.max_batch_size, self.network_summary[-1][2] ), dtype=np.float32)) 
  else:
   self.network_mem.append( gpuarray.empty((self.network_summary[-1][2], ), dtype=np.float32))

推断方法的实现

我们现在将在我们的SequentialNetwork类中添加两种推断方法——即，根据特定输入预测输出的方法。我们将首先称之为predict，这将由最终用户使用。在训练过程中，我们将不得不基于仅部分层的部分结果进行预测，我们将为此制作另一种方法，称为partial_predict。

让我们从实现predict开始。这将接受两个输入——以一维或二维 NumPy 数组形式的样本集，可能还有用户定义的 CUDA 流。我们将从样本（这里称为x）进行一些类型检查和格式化，记住样本将以行方式存储：

def predict(self, x, stream=None):

 if stream is None:
  stream = self.stream

 if type(x) != np.ndarray:
  temp = np.array(x, dtype = np.float32)
  x = temp

 if(x.size == self.network_mem[0].size):
  self.network_mem[0].set_async(x, stream=stream)
 else:

  if x.size > self.network_mem[0].size:
   raise Exception("Error: batch size too large for input.")

  x0 = np.zeros((self.network_mem[0].size,), dtype=np.float32)
  x0[0:x.size] = x.ravel()
  self.network_mem[0].set_async(x0.reshape( self.network_mem[0].shape), stream=stream)

 if(len(x.shape) == 2):
  batch_size = x.shape[0]
 else:
  batch_size = 1

现在，让我们执行实际的推断步骤。我们只需遍历整个神经网络，在每一层上执行eval_：

for i in xrange(len(self.network)):
 self.network[i].eval_(x=self.network_mem[i], y= self.network_mem[i+1], batch_size=batch_size, stream=stream)

现在，我们将提取 NN 的最终输出，GPU，并将其返回给用户。如果x中的样本数量实际上小于最大批处理大小，则在返回之前我们将适当地切片输出数组：

y = self.network_mem[-1].get_async(stream=stream)

if len(y.shape) == 2:
 y = y[0:batch_size, :]

return y

现在，完成了这一步，让我们实现partial_predict。让我们简要讨论一下这个想法。当我们处于训练过程中时，我们将评估一组样本，然后看看如何通过逐个添加delta到每个权重和偏差上会影响输出。为了节省时间，我们可以计算每一层的输出并将它们存储给定的样本集，然后只重新计算我们改变权重的层的输出，以及所有后续层的输出。我们很快就会更深入地了解这个想法，但现在，我们可以这样实现：

def partial_predict(self, layer_index=None, w_t=None, b_t=None, partial_mem=None, stream=None, batch_size=None, delta=None):

 self.network[layer_index].eval_(x=self.network_mem[layer_index], y = partial_mem[layer_index+1], batch_size=batch_size, stream = stream, w_t=w_t, b_t=b_t, delta=delta)

 for i in xrange(layer_index+1, len(self.network)):
  self.network[i].eval_(x=partial_mem[i], y =partial_mem[i+1], batch_size=batch_size, stream = stream)

梯度下降

我们现在将以批量随机梯度下降（BSGD）的形式对我们的 NN 进行训练方法的完整实现。让我们逐字逐句地思考这意味着什么。批量意味着这个训练算法将一次操作一组训练样本，而不是同时处理所有样本，而随机表示每个批次是随机选择的。梯度意味着我们将使用微积分中的梯度，这里是每个权重和偏差对损失函数的导数集合。最后，下降意味着我们试图减少损失函数——我们通过迭代地对权重和偏差进行微小的更改来实现这一点，通过减去梯度。

从微积分中我们知道，一个点的梯度总是指向最大增加的方向，其相反方向是最大减少的方向。因为我们想要减少，所以我们减去梯度。

我们现在将在我们的SequentialNetwork类中实现 BSGD 作为bsgd方法。让我们逐一讨论bsgd的输入参数：

• training将是一个二维 NumPy 数组的训练样本

• labels将是 NN 最终层的期望输出，对应于每个训练样本

• delta将指示我们在计算导数时应该增加权重多少

• max_streams将指示 BSGD 将在其上执行计算的最大并发 CUDA 流的数量

• batch_size将指示我们希望对每次更新权重计算损失函数的批次有多大

• epochs将指示我们对当前样本集的顺序进行多少次洗牌，分成一系列批次，然后进行 BSGD

• training_rate将指示我们使用梯度计算更新权重和偏差的速率

我们将像往常一样开始这个方法，并进行一些检查和类型转换，将 CUDA 流对象的集合设置为 Python 列表，并在另一个列表中分配一些额外需要的 GPU 内存：

def bsgd(self, training=None, labels=None, delta=None, max_streams = None, batch_size = None, epochs = 1, training_rate=0.01):

 training_rate = np.float32(training_rate)

 training = np.float32(training)
 labels = np.float32(labels)

 if( training.shape[0] != labels.shape[0] ):
  raise Exception("Number of training data points should be same as labels!")

 if max_streams is None:
  max_streams = self.max_streams

 if epochs is None:
 epochs = self.epochs

 if delta is None:
 delta = self.delta

 streams = []
 bgd_mem = []

 # create the streams needed for training
 for _ in xrange(max_streams):
  streams.append(drv.Stream())
  bgd_mem.append([])

 # allocate memory for each stream
 for i in xrange(len(bgd_mem)):
  for mem_bank in self.network_mem:
   bgd_mem[i].append( gpuarray.empty_like(mem_bank) )

现在，我们可以开始训练。我们将从每个epoch开始执行整个 BSGD 的迭代，对每个 epoch 对整个数据集进行随机洗牌。我们还将在终端打印一些信息，以便用户在训练过程中获得一些状态更新：

num_points = training.shape[0]

if batch_size is None:
 batch_size = self.max_batch_size

index = range(training.shape[0])

for k in xrange(epochs): 

 print '-----------------------------------------------------------'
 print 'Starting training epoch: %s' % k
 print 'Batch size: %s , Total number of training samples: %s' % (batch_size, num_points)
 print '-----------------------------------------------------------'

 all_grad = []

 np.random.shuffle(index)

现在，我们将循环遍历洗牌数据集中的每个批次。我们首先计算当前批次的熵，然后将其打印出来。如果用户看到熵的减少，那么他们将知道梯度下降在这里起作用：

for r in xrange(int(np.floor(training.shape[0]/batch_size))):

 batch_index = index[r*batch_size:(r+1)*batch_size] 

 batch_training = training[batch_index, :]
 batch_labels = labels[batch_index, :]

 batch_predictions = self.predict(batch_training)

 cur_entropy = cross_entropy(predictions=batch_predictions, ground_truth=batch_labels)

 print 'entropy: %s' % cur_entropy

我们现在将迭代我们的 NN 的每个密集层，计算整套权重和偏差的梯度。我们将把这些导数存储在扁平化（一维）数组中，这将对应于我们的 CUDA 内核中的w_t和b_t索引，它们也是扁平化的。由于我们将有多个流处理不同权重的不同输出，我们将使用 Python 队列容器来存储尚未处理的这一批权重和偏差：然后我们可以从该容器的顶部弹出值到下一个可用流（我们将这些存储为元组，第一个元素指示这是权重还是偏差）：

for i in xrange(len(self.network)):

 if self.network_summary[i][0] != 'dense':
  continue

 all_weights = Queue()

 grad_w = np.zeros((self.network[i].weights.size,), dtype=np.float32)
 grad_b = np.zeros((self.network[i].b.size,), dtype=np.float32)

 for w in xrange( self.network[i].weights.size ):
  all_weights.put( ('w', np.int32(w) ) )

 for b in xrange( self.network[i].b.size ):
  all_weights.put(('b', np.int32(b) ) )

现在，我们需要迭代每一个权重和偏差，我们可以使用while循环来检查我们刚刚设置的queue对象是否为空。我们将设置另一个队列stream_weights，这将帮助我们组织每个流处理的权重和偏差。适当设置权重和偏差输入后，我们现在可以使用partial_predict，使用当前流和相应的 GPU 内存数组：

请注意，我们已经对这批样本执行了predict来计算熵，所以我们现在可以对这批样本执行partial_predict，只要我们小心使用内存和层。

while not all_weights.empty():

 stream_weights = Queue()

 for j in xrange(max_streams):

  if all_weights.empty():
    break

  wb = all_weights.get()

  if wb[0] == 'w':
   w_t = wb[1]
   b_t = None
  elif wb[0] == 'b':
   b_t = wb[1]
   w_t = None

  stream_weights.put( wb )

  self.partial_predict(layer_index=i, w_t=w_t, b_t=b_t, partial_mem=bgd_mem[j], stream=streams[j], batch_size=batch_size, delta=delta)

我们只计算了一小部分权重和偏差的输出预测。我们将不得不为每个计算熵，然后将导数值存储在扁平化的数组中：

for j in xrange(max_streams):

 if stream_weights.empty():
  break

 wb = stream_weights.get()

 w_predictions = bgd_mem[j][-1].get_async(stream=streams[j])

 w_entropy = cross_entropy(predictions=w_predictions[ :batch_size,:], ground_truth=batch_labels)

 if wb[0] == 'w':
  w_t = wb[1]
  grad_w[w_t] = -(w_entropy - cur_entropy) / delta

 elif wb[0] == 'b':
  b_t = wb[1]
  grad_b[b_t] = -(w_entropy - cur_entropy) / delta

我们现在已经完成了while循环。一旦我们到达外部，我们将知道我们已经计算了这个特定层的所有权重和偏差的导数。在迭代到下一层之前，我们将把当前权重和偏差的梯度计算值附加到all_grad列表中。我们还将把扁平化的权重列表重新整形成原始形状：

all_grad.append([np.reshape(grad_w,self.network[i].weights.shape) , grad_b])

在迭代每一层之后，我们可以对这一批的 NN 的权重和偏差进行优化。请注意，如果training_rate变量远小于1，这将减少权重更新的速度：

for i in xrange(len(self.network)):
 if self.network_summary[i][0] == 'dense':
  new_weights = self.network[i].weights.get()
  new_weights += training_rate*all_grad[i][0]
  new_bias = self.network[i].b.get()
  new_bias += training_rate*all_grad[i][1]
  self.network[i].weights.set(new_weights)
  self.network[i].b.set(new_bias)

我们已经完全实现了（非常简单的）基于 GPU 的 DNN！

数据的调整和归一化

在我们继续训练和测试全新的 NN 之前，我们需要退后一步，谈谈数据调整和数据归一化。NN 对数值误差非常敏感，特别是当输入的规模差异很大时。这可以通过正确调整我们的训练数据来减轻，这意味着对于输入样本中的每个点，我们将计算所有样本中每个点的平均值和方差，然后在输入到 NN 进行训练或推断（预测）之前，对每个样本中的每个点减去平均值并除以标准差。这种方法称为归一化。让我们组合一个小的 Python 函数来为我们做这个：

def condition_data(data, means=None, stds=None):

 if means is None:
  means = np.mean(data, axis=0)

 if stds is None:
  stds = np.std(data, axis = 0)

 conditioned_data = data.copy()
 conditioned_data -= means
 conditioned_data /= stds

 return (conditioned_data, means, stds)

鸢尾花数据集

我们现在将为一个真实问题构建我们自己的 DNN：根据花瓣的测量来对花的类型进行分类。我们将使用众所周知的鸢尾花数据集。该数据集存储为逗号分隔值（CSV）文本文件，每行包含四个不同的数值（花瓣测量），然后是花的类型（这里有三个类别——山鸢尾、变色鸢尾和维吉尼亚鸢尾）。我们现在将设计一个小型 DNN，根据这个数据集对鸢尾花的类型进行分类。

在我们继续之前，请下载鸢尾花数据集并将其放入您的工作目录。这可以从 UC Irvine 机器学习存储库中获取，网址为：archive.ics.uci.edu/ml/machine-learning-databases/iris/iris.data。

我们将首先将此文件处理成适当的数据数组，以便用于训练和验证我们的 DNN。让我们从打开我们的主函数开始；我们需要将花的名称转换为 DNN 可以输出的实际类别，因此让我们创建一个小字典，为每个类别提供相应的标签。我们还将设置一些空列表来存储我们的训练数据和标签：

if __name__ == '__main__':
 to_class = { 'Iris-setosa' : [1,0,0] , 'Iris-versicolor' : [0,1,0], 'Iris-virginica' : [0,0,1]}

 iris_data = []
 iris_labels = []

现在，让我们从 CSV 文件中读取。我们将使用 Python 的csv模块中的reader函数，这是我们之前导入的：

with open('C:/Users/btuom/examples/9/iris.data', 'rb') as csvfile:
 csvreader = csv.reader(csvfile, delimiter=',')
 for row in csvreader:
  newrow = []
  if len(row) != 5:
   break
  for i in range(4):
   newrow.append(row[i])
  iris_data.append(newrow)
  iris_labels.append(to_class[row[4]])

我们现在将随机洗牌数据，并使用三分之二的样本作为训练数据。剩下的三分之一将用于测试（验证）数据：

iris_len = len(iris_data)
shuffled_index = list(range(iris_len))
np.random.shuffle(shuffled_index)

iris_data = np.float32(iris_data)
iris_labels = np.float32(iris_labels)
iris_data = iris_data[shuffled_index, :]
iris_labels = iris_labels[shuffled_index,:]

t_len = (2*iris_len) // 3

iris_train = iris_data[:t_len, :]
label_train = iris_labels[:t_len, :]

iris_test = iris_data[t_len:,:]
label_test = iris_labels[t_len:, :]

现在，最后，我们可以开始构建我们的 DNN！首先，让我们创建一个SequentialNetwork对象。我们将max_batch_size设置为32：

sn = SequentialNetwork( max_batch_size=32 )

现在，让我们创建我们的 NN。这将包括四个密集层（两个隐藏层）和一个 softmax 层。我们将逐层增加神经元的数量，直到最后一层，最后一层只有三个输出（每个类别一个）。每层神经元数量的增加允许我们捕捉数据的一些细微差别：

sn.add_layer({'type' : 'dense', 'num_inputs' : 4, 'num_outputs' : 10, 'relu': True, 'sigmoid': False, 'weights' : None, 'bias' : None} ) 
sn.add_layer({'type' : 'dense', 'num_inputs' : 10, 'num_outputs' : 15, 'relu': True, 'sigmoid': False, 'weights': None, 'bias' : None} ) 
sn.add_layer({'type' : 'dense', 'num_inputs' : 15, 'num_outputs' : 20, 'relu': True, 'sigmoid': False, 'weights': None, 'bias' : None} ) 
sn.add_layer({'type' : 'dense', 'num_inputs' : 20, 'num_outputs' : 3, 'relu': True, 'sigmoid': False, 'weights': None , 'bias': None } ) 
sn.add_layer({'type' : 'softmax'})

我们现在将调整我们的训练数据，并使用我们刚刚实现的 BSGD 方法进行训练。我们将使用batch_size设置为16，max_streams设置为10，epochs的数量设置为 100，delta设置为 0.0001，training_rate设置为 1——这些将是几乎任何现代 GPU 可接受的参数。我们还将计时训练过程，这可能会非常耗时。

ctrain, means, stds = condition_data(iris_train)

t1 = time()
sn.bsgd(training=ctrain, labels=label_train, batch_size=16, max_streams=20, epochs=100 , delta=0.0001, training_rate=1)
training_time = time() - t1

现在，我们的 DNN 已经完全训练好了。我们准备开始验证过程！让我们设置一个名为hits的 Python 变量来计算正确分类的总数。我们还需要对验证/测试数据进行条件设置。还有一件事——我们通过 DNN 的 softmax 层的最大值对应的索引来确定类别。我们可以使用 NumPy 的argmax函数来检查这是否给出了正确的分类，就像这样：

hits = 0
ctest, _, _ = condition_data(iris_test, means=means, stds=stds)
for i in range(ctest.shape[0]):
 if np.argmax(sn.predict(ctest[i,:])) == np.argmax( label_test[i,:]):
  hits += 1

现在，我们准备检查我们的 DNN 实际工作得有多好。让我们输出准确率以及总训练时间：

print 'Percentage Correct Classifications: %s' % (float(hits ) / ctest.shape[0])
print 'Total Training Time: %s' % training_time

现在，我们完成了。我们现在可以完全用 Python 和 CUDA 实现一个 DNN！一般来说，您可以期望在这个特定问题上的准确率在 80%-97%之间，使用任何 Pascal 级别的 GPU 的训练时间为 10-20 分钟。

本章的代码可在本书的 GitHub 存储库的适当目录下的deep_neural_network.py文件中找到。

总结

在本章中，我们首先给出了人工神经网络的定义，并向您展示了如何将单个 AN 组合成密集层，然后再将其组合成完整的深度神经网络。然后，我们在 CUDA-C 中实现了一个密集层，并制作了一个相应的 Python 包装类。我们还包括了在密集层的输出上添加 ReLU 和 sigmoid 层的功能。我们看到了使用 softmax 层的定义和动机，这用于分类问题，然后在 CUDA-C 和 Python 中实现了这一功能。最后，我们实现了一个 Python 类，以便我们可以从先前的类构建一个顺序前馈 DNN；我们实现了一个交叉熵损失函数，然后在我们的梯度下降实现中使用这个损失函数来训练我们 DNN 中的权重和偏差。最后，我们使用我们的实现在真实数据集上构建、训练和测试了一个 DNN。

现在我们对我们的 CUDA 编程能力有了很大的自信，因为我们可以编写自己基于 GPU 的 DNN！我们现在将在接下来的两章中学习一些非常高级的内容，我们将看看如何编写我们自己的接口到编译后的 CUDA 代码，以及一些关于 NVIDIA GPU 非常技术性的细节。

问题

1. 假设您构建了一个 DNN，并在训练后，它只产生垃圾。经过检查，您发现所有的权重和偏差要么是巨大的数字，要么是 NaN。问题可能是什么？

2. 小training_rate值可能存在的一个问题是什么？

3. 大training_rate值可能存在的一个问题是什么？

4. 假设我们想要训练一个 DNN，将多个标签分配给动物的图像（“黏滑的”，“毛茸茸的”，“红色的”，“棕色的”等等）。在 DNN 的末尾应该使用 sigmoid 还是 softmax 层？

5. 假设我们想要将单个动物的图像分类为猫或狗。我们应该使用 sigmoid 还是 softmax？

6. 如果我们减小批量大小，梯度下降训练过程中权重和偏差的更新会更多还是更少？

第十章：使用已编译的 GPU 代码

在本书的过程中，我们通常依赖 PyCUDA 库自动为我们接口我们的内联 CUDA-C 代码，使用即时编译和与 Python 代码的链接。然而，我们可能还记得，有时编译过程可能需要一段时间。在第三章中，使用 PyCUDA 入门，我们甚至详细看到编译过程如何导致减速，以及内联代码何时被编译和保留可能是相当随意的。在某些情况下，这可能会给应用程序带来不便，或者在实时系统的情况下甚至是不可接受的。

为此，我们最终将看到如何从 Python 使用预编译的 GPU 代码。特别是，我们将看看三种不同的方法来做到这一点。首先，我们将看看如何通过编写一个主机端 CUDA-C 函数来间接启动 CUDA 内核。这种方法将涉及使用标准 Python Ctypes 库调用主机端函数。其次，我们将把我们的内核编译成所谓的 PTX 模块，这实际上是一个包含已编译二进制 GPU 的 DLL 文件。然后，我们可以使用 PyCUDA 加载此文件并直接启动我们的内核。最后，我们将通过查看如何编写我们自己的完整 Ctypes 接口来结束本章，以使用 CUDA Driver API。然后，我们可以使用 Driver API 中的适当函数加载我们的 PTX 文件并启动内核。

本章的学习成果如下：

• 使用 Ctypes 模块启动编译后（主机端）的代码

• 使用 Ctypes 使用主机端 CUDA C 包装器从 Python 启动内核

• 如何将 CUDA C 模块编译为 PTX 文件

• 如何将 PTX 模块加载到 PyCUDA 中以启动预编译的内核

• 如何编写自定义 Python 接口以使用 CUDA Driver API

使用 Ctypes 启动编译后的代码

我们现在将简要概述 Python 标准库中的 Ctypes 模块。Ctypes 用于调用来自 Linux.so（共享对象）或 Windows.DLL（动态链接库）预编译二进制文件的函数。这将使我们摆脱纯 Python 的世界，并与已用编译语言编写的库和代码进行接口，特别是 C 和 C++ - 恰好 Nvidia 只为与我们的 CUDA 设备进行接口提供这样的预编译二进制文件，因此如果我们想绕过 PyCUDA，我们将不得不使用 Ctypes。

让我们从一个非常基本的例子开始：我们将向您展示如何直接从 Ctypes 调用printf。打开一个 IPython 实例，键入import ctypes。现在我们将看看如何从 Ctypes 调用标准的printf函数。首先，我们必须导入适当的库：在 Linux 中，通过键入libc = ctypes.CDLL('libc.so.6')加载 LibC 库（在 Windows 中，将'libc.so.6'替换为'msvcrt.dll'）。现在我们可以通过在 IPython 提示符中键入libc.printf("Hello from ctypes!\n")直接调用printf。自己试试吧！

现在让我们试试其他东西：从 IPython 键入libc.printf("Pi is approximately %f.\n", 3.14)；您应该会收到一个错误。这是因为3.14没有适当地从 Python 浮点变量转换为 C 双精度变量 - 我们可以使用 Ctypes 这样做：

libc.printf("Pi is approximately %f.\n", ctypes.c_double(3.14)) 

输出应该如预期那样。与从 PyCUDA 启动 CUDA 内核的情况一样，我们必须同样小心地将输入转换为 Ctypes 函数。

请务必确保将任何从 Python 使用 Ctypes 调用的函数的输入适当地转换为适当的 C 数据类型（在 Ctypes 中，这些类型以 c_ 开头：c_float、c_double、c_char、c_int等）。

再次重温 Mandelbrot 集

让我们重新审视一下我们在第一章和第三章中看到的 Mandelbrot 集合，为什么使用 GPU 编程？和使用 PyCUDA 入门。首先，我们将编写一个完整的 CUDA 核函数，它将根据一组特定的参数计算 Mandelbrot 集合，以及一个适当的主机端包装函数，我们稍后可以从 Ctypes 接口调用。我们将首先将这些函数编写到一个单独的 CUDA-C.cu源文件中，然后使用 NVCC 编译成 DLL 或.so二进制文件。最后，我们将编写一些 Python 代码，以便我们可以运行我们的二进制代码并显示 Mandelbrot 集合。

我们现在将运用我们对 Ctypes 的知识，从 Python 中启动一个预编译的 CUDA 核函数，而不需要 PyCUDA 的任何帮助。这将要求我们在 CUDA-C 中编写一个主机端核函数启动器包装函数，我们可以直接调用，它本身已经编译成了一个动态库二进制文件，其中包含任何必要的 GPU 代码——即在 Windows 上的动态链接库（DLL）二进制文件，或者在 Linux 上的共享对象（so）二进制文件。

当然，我们将首先编写我们的 CUDA-C 代码，所以打开你最喜欢的文本编辑器并跟着做。我们将从标准的include语句开始：

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

我们现在将直接开始编写我们的核函数。请注意代码中的extern "C"，这将允许我们在外部链接到这个函数：

extern "C" __global__ void mandelbrot_ker(float * lattice, float * mandelbrot_graph, int max_iters, float upper_bound_squared, int lattice_size)
{

让我们思考一分钟关于这将如何工作：我们将使用一个单一的一维数组来存储实部和虚部，称为lattice，其长度为lattice_size。我们将使用这个数组来计算一个形状为(lattice_size, lattice_size)的二维 Mandelbrot 图形，存储在预先分配的数组mandelbrot_graph中。我们将指定每个点检查发散的迭代次数为max_iters，通过使用upper_bound_squared提供其平方值来指定之前的最大上限值。（我们稍后会看一下使用平方的动机。）

我们将在一维网格/块结构上启动这个核函数，每个线程对应于 Mandelbrot 集合图像中的一个点。然后我们可以确定相应点的实部/虚部 lattice 值，如下所示：

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

    if ( tid < lattice_size*lattice_size )
    {
        int i = tid % lattice_size;
        int j = lattice_size - 1 - (tid / lattice_size);

        float c_re = lattice[i];
        float c_im = lattice[j];

让我们花一分钟来谈谈这个。首先，记住我们可能需要使用稍多于必要的线程，所以重要的是我们检查线程 ID 是否对应于输出图像中的某个点，使用if语句。让我们还记住，输出数组mandelbrot_graph将以按行方式存储为一维数组，表示为二维图像，我们将使用tid作为写入该数组的索引。我们将使用i和j，以及复平面上图形的x和y坐标。由于 lattice 是一系列从小到大排序的实值，我们将不得不颠倒它们的顺序以获得适当的虚值。另外，请注意，我们将在这里使用普通的浮点数，而不是某种结构或对象来表示复数值。由于每个复数中都有实部和虚部，我们将在这里使用两个浮点数来存储与该线程的 lattice 点对应的复数（c_re和c_im）。

我们将设置两个更多的变量来处理发散检查，z_re和z_im，并在检查发散之前将该线程的图上的点的初始值设置为1：

        float z_re = 0.0f;
        float z_im = 0.0f;

        mandelbrot_graph[tid] = 1;

现在我们将检查发散；如果在max_iters次迭代后发散，我们将把点设置为0。否则，它将保持为 1：

        for (int k = 0; k < max_iters; k++)
        {
            float temp;

            temp = z_re*z_re - z_im*z_im + c_re;
            z_im = 2*z_re*z_im + c_im;
            z_re = temp;

            if ( (z_re*z_re + z_im*z_im) > upper_bound_squared )
            {
                mandelbrot_graph[tid] = 0;
                break;
            }
        }

在我们继续之前，让我们谈一分钟关于这一块代码。让我们记住，曼德勃罗集的每次迭代都是通过复数乘法和加法来计算的，例如，z_new = z*z + c。由于我们不是使用将为我们处理复数值的类，前面的操作正是我们需要做的，以计算z的新的实部和虚部值。我们还需要计算绝对值并查看是否超过特定值——记住，复数的绝对值，c = x + iy，是用√(x²+y²)来计算的。在这里计算上限的平方然后将其插入内核中，实际上会节省我们一些时间，因为这样可以节省我们在每次迭代中计算z_re*z_re + z_im*z_im的平方根的时间。

我们现在基本上已经完成了这个内核——我们只需要关闭if语句并从内核返回，然后我们就完成了：

    }
    return;
}

然而，我们还没有完全完成。我们需要编写一个只有extern "C"的主机端包装函数，在 Linux 的情况下，以及在 Windows 的情况下，只有extern "C" __declspec(dllexport)。 （与编译的 CUDA 内核相反，如果我们想要能够从 Ctypes 在 Windows 中访问主机端函数，这个额外的单词是必要的。）我们放入这个函数的参数将直接对应于进入内核的参数，除了这些参数将存储在主机上：

extern "C" __declspec(dllexport) void launch_mandelbrot(float * lattice,  float * mandelbrot_graph, int max_iters, float upper_bound, int lattice_size)
{

现在，我们将需要分配足够的内存来存储在 GPU 上的晶格和输出，然后使用cudaMemcpy将晶格复制到 GPU 上：

    int num_bytes_lattice = sizeof(float) * lattice_size;
    int num_bytes_graph = sizeof(float)* lattice_size*lattice_size;

    float * d_lattice;
    float * d_mandelbrot_graph;

    cudaMalloc((float **) &d_lattice, num_bytes_lattice);
    cudaMalloc((float **) &d_mandelbrot_graph, num_bytes_graph);

    cudaMemcpy(d_lattice, lattice, num_bytes_lattice, cudaMemcpyHostToDevice);

像我们的许多其他内核一样，我们将在一维网格上启动大小为 32 的一维块。我们将取输出点数除以 32 的上限值，以确定网格大小，如下所示：

    int grid_size = (int)  ceil(  ( (double) lattice_size*lattice_size ) / ( (double) 32 ) );

现在我们准备使用传统的 CUDA-C 三角形括号来启动我们的内核，以指定网格和块大小。请注意，我们在这里提前求出了上限的平方：

    mandelbrot_ker <<< grid_size, 32 >>> (d_lattice,  d_mandelbrot_graph, max_iters, upper_bound*upper_bound, lattice_size);

现在我们只需要在完成这些操作后将输出复制到主机上，然后在适当的数组上调用cudaFree。然后我们可以从这个函数返回：

    cudaMemcpy(mandelbrot_graph, d_mandelbrot_graph, num_bytes_graph, cudaMemcpyDeviceToHost);    
    cudaFree(d_lattice);
    cudaFree(d_mandelbrot_graph);
}

有了这个，我们已经完成了所有需要的 CUDA-C 代码。将其保存到名为mandelbrot.cu的文件中，然后继续下一步。

您还可以从github.com/btuomanen/handsongpuprogramming/blob/master/10/mandelbrot.cu下载此文件。

编译代码并与 Ctypes 进行接口

现在让我们将刚刚编写的代码编译成 DLL 或.so二进制文件。这实际上相当简单：如果你是 Linux 用户，请在命令行中输入以下内容将此文件编译成mandelbrot.so：

nvcc -Xcompiler -fPIC -shared -o mandelbrot.so mandelbrot.cu

如果你是 Windows 用户，请在命令行中输入以下内容将文件编译成mandelbrot.dll：

nvcc -shared -o mandelbrot.dll mandelbrot.cu

现在我们可以编写我们的 Python 接口。我们将从适当的导入语句开始，完全排除 PyCUDA，只使用 Ctypes。为了方便使用，我们将直接从 Ctypes 导入所有类和函数到默认的 Python 命名空间中，如下所示：

from __future__ import division
from time import time
import matplotlib
from matplotlib import pyplot as plt
import numpy as np
from ctypes import *

让我们使用 Ctypes 为launch_mandelbrot主机端函数设置一个接口。首先，我们将不得不加载我们编译的 DLL 或.so文件，Linux 用户当然需要将文件名更改为mandelbrot.so：

mandel_dll = CDLL('./mandelbrot.dll')

现在我们可以从库中获取对launch_mandelbrot的引用，就像这样；我们将简称它为mandel_c：

mandel_c = mandel_dll.launch_mandelbrot

现在，在使用 Ctypes 调用函数之前，我们将不得不让 Ctypes 知道输入类型是什么。让我们记住，对于launch_mandelbrot，输入是float-pointer，float-pointer，integer，float和integer。我们使用argtypes参数设置这一点，使用适当的 Ctypes 数据类型（c_float，c_int），以及 Ctypes 的POINTER类：

mandel_c.argtypes = [POINTER(c_float), POINTER(c_float), c_int, c_float, c_int]

现在让我们编写一个 Python 函数来为我们运行这个。我们将使用breadth指定正方形输出图像的宽度和高度，并且在复杂格的实部和虚部中指定最小和最大值。我们还将指定最大迭代次数以及上限：

def mandelbrot(breadth, low, high, max_iters, upper_bound):

现在，我们将使用 NumPy 的linspace函数创建我们的格点数组，就像这样：

 lattice = np.linspace(low, high, breadth, dtype=np.float32)

让我们记住，我们将不得不传递一个预先分配的浮点数组给launch_mandelbrot，以便以输出图的形式得到输出。我们可以通过调用 NumPy 的empty命令来设置一个适当形状和大小的数组，这将充当 C 的malloc调用：

    out = np.empty(shape=(lattice.size,lattice.size), dtype=np.float32)

现在，我们准备计算 Mandelbrot 图。请注意，我们可以通过使用它们的ctypes.data_as方法和相应的类型将 NumPy 数组传递给 C。在我们这样做之后，我们可以返回输出；也就是说，Mandelbrot 图以二维 NumPy 数组的形式：

 mandel_c(lattice.ctypes.data_as(POINTER(c_float)), out.ctypes.data_as(POINTER(c_float)), c_int(max_iters), c_float(upper_bound), c_int(lattice.size) ) 
 return out

现在，让我们编写我们的主函数来计算、计时和使用 Matplotlib 查看 Mandelbrot 图：

if __name__ == '__main__':
    t1 = time()
    mandel = mandelbrot(512,-2,2,256, 2)
    t2 = time()
    mandel_time = t2 - t1
    print 'It took %s seconds to calculate the Mandelbrot graph.' % mandel_time
    plt.figure(1)
    plt.imshow(mandel, extent=(-2, 2, -2, 2))
    plt.show()

我们现在将尝试运行这个。您应该会得到一个看起来与第一章的 Mandelbrot 图以及第三章的为什么使用 GPU 编程和使用 PyCUDA 入门中的图形完全相同的输出：

这个 Python 示例的代码也可以在 GitHub 存储库的mandelbrot_ctypes.py文件中找到。

编译和启动纯 PTX 代码

我们刚刚看到了如何从 Ctypes 调用纯 C 函数。在某些方面，这可能看起来有点不够优雅，因为我们的二进制文件必须包含主机代码以及编译后的 GPU 代码，这可能看起来很麻烦。我们是否可以只使用纯粹的编译后的 GPU 代码，然后适当地将其启动到 GPU 上，而不是每次都编写一个 C 包装器？幸运的是，我们可以。

NVCC 编译器将 CUDA-C 编译为PTX（Parallel Thread Execution），这是一种解释的伪汇编语言，与 NVIDIA 的各种 GPU 架构兼容。每当您使用 NVCC 将使用 CUDA 核心的程序编译为可执行的 EXE、DLL、.so或 ELF 文件时，该文件中将包含该核心的 PTX 代码。我们还可以直接编译具有 PTX 扩展名的文件，其中将仅包含从编译后的 CUDA .cu 文件中编译的 GPU 核心。幸运的是，PyCUDA 包括一个接口，可以直接从 PTX 加载 CUDA 核心，使我们摆脱了即时编译的枷锁，同时仍然可以使用 PyCUDA 的所有其他好功能。

现在让我们将刚刚编写的 Mandelbrot 代码编译成一个 PTX 文件；我们不需要对它进行任何更改。只需在 Linux 或 Windows 的命令行中输入以下内容：

nvcc -ptx -o mandelbrot.ptx mandelbrot.cu

现在让我们修改上一节的 Python 程序，以使用 PTX 代码。我们将从导入中删除ctypes并添加适当的 PyCUDA 导入：

from __future__ import division
from time import time
import matplotlib
from matplotlib import pyplot as plt
import numpy as np
import pycuda
from pycuda import gpuarray
import pycuda.autoinit

现在让我们使用 PyCUDA 的module_from_file函数加载 PTX 文件，就像这样：

mandel_mod = pycuda.driver.module_from_file('./mandelbrot.ptx')

现在我们可以使用get_function来获取对我们的核心的引用，就像我们用 PyCUDA 的SourceModule一样：

mandel_ker = mandel_mod.get_function('mandelbrot_ker')

我们现在可以重写 Mandelbrot 函数，以处理使用适当的gpuarray对象和typecast输入的核心。（我们不会逐行讨论这个，因为在这一点上它的功能应该是显而易见的。）：

def mandelbrot(breadth, low, high, max_iters, upper_bound):
    lattice = gpuarray.to_gpu(np.linspace(low, high, breadth, dtype=np.   
    out_gpu = gpuarray.empty(shape=(lattice.size,lattice.size), dtype=np.float32)
    gridsize = int(np.ceil(lattice.size**2 / 32))
    mandel_ker(lattice, out_gpu, np.int32(256), np.float32(upper_bound**2), np.int32(lattice.size), grid=(gridsize, 1, 1), block=(32,1,1))
    out = out_gpu.get()

    return out

main函数将与上一节完全相同：

if __name__ == '__main__':
    t1 = time()
    mandel = mandelbrot(512,-2,2,256,2)
    t2 = time()
    mandel_time = t2 - t1
    print 'It took %s seconds to calculate the Mandelbrot graph.' % mandel_time
    plt.figure(1)
    plt.imshow(mandel, extent=(-2, 2, -2, 2))
    plt.show()

现在，尝试运行这个来确保输出是正确的。您可能还会注意到与 Ctypes 版本相比的一些速度改进。

这段代码也可以在本书 GitHub 存储库的“10”目录下的mandelbrot_ptx.py文件中找到。

为 CUDA Driver API 编写包装器

现在我们将看看如何使用 Ctypes 编写我们自己的包装器，用于一些预打包的二进制 CUDA 库函数。特别是，我们将为 CUDA 驱动程序 API 编写包装器，这将允许我们执行所有基本 GPU 使用所需的操作，包括 GPU 初始化、内存分配/传输/释放、内核启动和上下文创建/同步/销毁。这是一个非常强大的知识；它将允许我们在不经过 PyCUDA 的情况下使用 GPU，也不需要编写任何繁琐的主机端 C 函数包装器。

现在我们将编写一个小模块，它将作为CUDA 驱动程序 API的包装库。让我们谈一分钟这对我们意味着什么。驱动程序 API 与CUDA Runtime API略有不同，技术性稍高，后者是我们在 CUDA-C 文本中一直在使用的。驱动程序 API 旨在与常规 C/C++编译器一起使用，而不是与 NVCC 一起使用，具有一些不同的约定，如使用cuLaunchKernel函数启动内核，而不是使用<<< gridsize, blocksize >>>括号表示法。这将允许我们直接访问使用 Ctypes 从 PTX 文件启动内核所需的函数。

让我们通过将所有 Ctypes 导入模块的命名空间，并导入 sys 模块来开始编写此模块。我们将通过使用sys.platform检查系统的操作系统（nvcuda.dll或libcuda.so）来加载适当的库文件，使我们的模块可以在 Windows 和 Linux 上使用。

from ctypes import *
import sys
if 'linux' in sys.platform:
 cuda = CDLL('libcuda.so')
elif 'win' in sys.platform:
 cuda = CDLL('nvcuda.dll')

我们已经成功加载了 CUDA 驱动程序 API，现在我们可以开始为基本 GPU 使用编写必要函数的包装器。随着我们的进行，我们将查看每个驱动程序 API 函数的原型，这通常是在编写 Ctypes 包装器时必要的。

鼓励读者在官方 Nvidia CUDA 驱动程序 API 文档中查找我们将在本节中使用的所有函数，该文档可在此处找到：docs.nvidia.com/cuda/cuda-driver-api/。

让我们从驱动程序 API 中最基本的函数cuInit开始，它将初始化驱动程序 API。这需要一个用于标志的无符号整数作为输入参数，并返回类型为 CUresult 的值，实际上只是一个整数值。我们可以这样编写我们的包装器：

cuInit = cuda.cuInit
cuInit.argtypes = [c_uint]
cuInit.restype = int

现在让我们开始下一个函数cuDeviceCount，它将告诉我们在计算机上安装了多少个 NVIDIA GPU。它以整数指针作为其唯一输入，实际上是通过引用返回的单个整数输出值。返回值是另一个 CUresult 整数——所有函数都将使用 CUresult，这是所有驱动程序 API 函数的错误值的标准化。例如，如果我们看到任何函数返回0，这意味着结果是CUDA_SUCCESS，而非零结果将始终意味着错误或警告：

cuDeviceGetCount = cuda.cuDeviceGetCount
cuDeviceGetCount.argtypes = [POINTER(c_int)]
cuDeviceGetCount.restype = int

现在让我们为cuDeviceGet编写一个包装器，它将通过引用在第一个输入中返回设备句柄。这将对应于第二个输入中给定的序号 GPU。第一个参数的类型是CUdevice *，实际上只是一个整数指针：

cuDeviceGet = cuda.cuDeviceGet
cuDeviceGet.argtypes = [POINTER(c_int), c_int]
cuDeviceGet.restype = int

让我们记住，每个 CUDA 会话都需要至少一个 CUDA 上下文，可以将其类比为在 CPU 上运行的进程。由于 Runtime API 会自动处理这一点，在这里我们将不得不在使用设备之前手动在设备上创建上下文（使用设备句柄），并且在 CUDA 会话结束时销毁此上下文。

我们可以使用cuCtxCreate函数创建一个 CUDA 上下文，它当然会创建一个上下文。让我们看看文档中列出的原型：

 CUresult cuCtxCreate ( CUcontext* pctx, unsigned int flags, CUdevice dev )

当然，返回值是CUresult。第一个输入是指向名为CUcontext的类型的指针，实际上它本身是 CUDA 内部使用的特定 C 结构的指针。由于我们从 Python 对CUcontext的唯一交互将是保持其值以在其他函数之间传递，我们可以将CUcontext存储为 C void *类型，用于存储任何类型的通用指针地址。由于这实际上是指向 CU 上下文的指针（再次，它本身是指向内部数据结构的指针——这是另一个按引用返回值），我们可以将类型设置为普通的void *，这是 Ctypes 中的c_void_p类型。第二个值是无符号整数，而最后一个值是要在其上创建新上下文的设备句柄——让我们记住这实际上只是一个整数。我们现在准备为cuCtxCreate创建包装器：

cuCtxCreate = cuda.cuCtxCreate
cuCtxCreate.argtypes = [c_void_p, c_uint, c_int]
cuCtxCreate.restype = int

您可以始终在 C/C++（Ctypes 中的c_void_p）中使用void *类型指向任意数据或变量，甚至结构和对象，其定义可能不可用。

下一个函数是cuModuleLoad，它将为我们加载一个 PTX 模块文件。第一个参数是一个 CUmodule 的引用（同样，我们可以在这里使用c_void_p），第二个是文件名，这将是一个典型的以空字符结尾的 C 字符串——这是一个char *，或者在 Ctypes 中是c_char_p：

cuModuleLoad = cuda.cuModuleLoad
cuModuleLoad.argtypes = [c_void_p, c_char_p]
cuModuleLoad.restype = int

下一个函数用于同步当前 CUDA 上下文中的所有启动操作，并称为cuCtxSynchronize（不带参数）：

cuCtxSynchronize = cuda.cuCtxSynchronize
cuCtxSynchronize.argtypes = []
cuCtxSynchronize.restype = int

下一个函数用于从加载的模块中检索内核函数句柄，以便我们可以将其启动到 GPU 上，这与 PyCUDA 的get_function方法完全对应，这一点我们已经看过很多次了。文档告诉我们原型是CUresult cuModuleGetFunction ( CUfunction* hfunc, CUmodule hmod, const char* name )。现在我们可以编写包装器：

cuModuleGetFunction = cuda.cuModuleGetFunction
 cuModuleGetFunction.argtypes = [c_void_p, c_void_p, c_char_p ]
 cuModuleGetFunction.restype = int

现在让我们为标准动态内存操作编写包装器；这些将是必要的，因为我们将不再有使用 PyCUDA gpuarray 对象的虚荣。这些实际上与我们之前使用过的 CUDA 运行时操作几乎相同；也就是说，cudaMalloc，cudaMemcpy和cudaFree：

cuMemAlloc = cuda.cuMemAlloc
cuMemAlloc.argtypes = [c_void_p, c_size_t]
cuMemAlloc.restype = int

cuMemcpyHtoD = cuda.cuMemcpyHtoD
cuMemcpyHtoD.argtypes = [c_void_p, c_void_p, c_size_t]
cuMemAlloc.restype = int

cuMemcpyDtoH = cuda.cuMemcpyDtoH
cuMemcpyDtoH.argtypes = [c_void_p, c_void_p, c_size_t]
cuMemcpyDtoH.restype = int

cuMemFree = cuda.cuMemFree
cuMemFree.argtypes = [c_void_p] 
cuMemFree.restype = int

现在，我们将为cuLaunchKernel函数编写一个包装器。当然，这是我们将用来在 GPU 上启动 CUDA 内核的函数，前提是我们已经初始化了 CUDA Driver API，设置了上下文，加载了一个模块，分配了内存并配置了输入，并且已经从加载的模块中提取了内核函数句柄。这个函数比其他函数复杂一些，所以我们将看一下原型：

CUresult cuLaunchKernel ( CUfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, CUstream hStream, void** kernelParams, void** extra )  

第一个参数是我们要启动的内核函数的句柄，我们可以表示为c_void_p。六个gridDim和blockDim参数用于指示网格和块的维度。无符号整数sharedMemBytes用于指示在内核启动时为每个块分配多少字节的共享内存。CUstream hStream是一个可选参数，我们可以使用它来设置自定义流，或者如果希望使用默认流，则设置为 NULL（0），我们可以在 Ctypes 中表示为c_void_p。最后，kernelParams和extra参数用于设置内核的输入；这些有点复杂，所以现在只需知道我们也可以将这些表示为c_void_p：

cuLaunchKernel = cuda.cuLaunchKernel
cuLaunchKernel.argtypes = [c_void_p, c_uint, c_uint, c_uint, c_uint, c_uint, c_uint, c_uint, c_void_p, c_void_p, c_void_p]
cuLaunchKernel.restype = int

现在我们还有最后一个函数要为cuCtxDestroy编写一个包装器。我们在 CUDA 会话结束时使用它来销毁 GPU 上的上下文。唯一的输入是一个CUcontext对象，由c_void_p表示：

cuCtxDestroy = cuda.cuCtxDestroy
cuCtxDestroy.argtypes = [c_void_p]
cuCtxDestroy.restype = int

让我们把这个保存到cuda_driver.py文件中。我们现在已经完成了 Driver API 包装器模块！接下来，我们将看看如何仅使用我们的模块和 Mandelbrot PTX 加载一个 PTX 模块并启动一个内核。

这个示例也可以在本书的 GitHub 存储库中的cuda_driver.py文件中找到。

使用 CUDA Driver API

我们现在将翻译我们的小曼德布洛特生成程序，以便我们可以使用我们的包装库。让我们从适当的导入语句开始；注意我们如何将所有的包装器加载到当前命名空间中：

from __future__ import division
from time import time
import matplotlib
from matplotlib import pyplot as plt
import numpy as np
from cuda_driver import *

让我们把所有的 GPU 代码放到mandelbrot函数中，就像以前一样。我们将从使用cuInit初始化 CUDA Driver API 开始，然后检查系统上是否安装了至少一个 GPU，否则会引发异常：

def mandelbrot(breadth, low, high, max_iters, upper_bound):
 cuInit(0)
 cnt = c_int(0)
 cuDeviceGetCount(byref(cnt))
 if cnt.value == 0:
  raise Exception('No GPU device found!')

注意这里的byref：这是 Ctypes 中引用操作符(&)的等价物。我们现在将再次应用这个想法，记住设备句柄和 CUDA 上下文可以用 Ctypes 表示为c_int和c_void_p：

 cuDevice = c_int(0)
 cuDeviceGet(byref(cuDevice), 0)
 cuContext = c_void_p()
 cuCtxCreate(byref(cuContext), 0, cuDevice)

我们现在将加载我们的 PTX 模块，记得要使用c_char_p将文件名转换为 C 字符串：

 cuModule = c_void_p()
 cuModuleLoad(byref(cuModule), c_char_p('./mandelbrot.ptx'))

现在我们将在主机端设置晶格，以及一个名为graph的用于在主机端存储输出的全零 NumPy 数组。我们还将为晶格和图形输出在 GPU 上分配内存，然后使用cuMemcpyHtoD将晶格复制到 GPU 上：

 lattice = np.linspace(low, high, breadth, dtype=np.float32)
 lattice_c = lattice.ctypes.data_as(POINTER(c_float))
 lattice_gpu = c_void_p(0)
 graph = np.zeros(shape=(lattice.size, lattice.size), dtype=np.float32)
 cuMemAlloc(byref(lattice_gpu), c_size_t(lattice.size*sizeof(c_float)))
 graph_gpu = c_void_p(0)
 cuMemAlloc(byref(graph_gpu), c_size_t(lattice.size**2 * sizeof(c_float)))
 cuMemcpyHtoD(lattice_gpu, lattice_c, c_size_t(lattice.size*sizeof(c_float)))

现在我们将使用cuModuleGetFunction获取 Mandelbrot 内核的句柄，并设置一些输入：

 mandel_ker = c_void_p(0)
 cuModuleGetFunction(byref(mandel_ker), cuModule, c_char_p('mandelbrot_ker'))
 max_iters = c_int(max_iters)
 upper_bound_squared = c_float(upper_bound**2)
 lattice_size = c_int(lattice.size)

下一步有点复杂。在继续之前，我们必须了解参数是如何通过cuLaunchKernel传递到 CUDA 内核中的。让我们先看看 CUDA-C 中是如何工作的。

我们将输入参数在kernelParams中表达为一个void *值的数组，它们本身是指向我们希望插入内核的输入的指针。对于我们的曼德布洛特内核，它看起来像这样：

void * mandel_params [] = {&lattice_gpu, &graph_gpu, &max_iters, &upper_bound_squared, &lattice_size};

现在让我们看看如何在 Ctypes 中表达这一点，这并不是立即显而易见的。首先，让我们将所有的输入放入一个 Python 列表中，按正确的顺序：

mandel_args0 = [lattice_gpu, graph_gpu, max_iters, upper_bound_squared, lattice_size ]

现在我们需要每个值的指针，将其类型转换为void *类型。让我们使用 Ctypes 函数addressof来获取每个 Ctypes 变量的地址（类似于byref，只是不绑定到特定类型），然后将其转换为c_void_p。我们将这些值存储在另一个列表中：

mandel_args = [c_void_p(addressof(x)) for x in mandel_args0]

现在让我们使用 Ctypes 将这个 Python 列表转换成一个void *指针数组，就像这样：

 mandel_params = (c_void_p * len(mandel_args))(*mandel_args)

现在我们可以设置网格的大小，就像以前一样，并使用cuLaunchKernel启动我们的内核，使用这组参数。然后我们在之后同步上下文：

 gridsize = int(np.ceil(lattice.size**2 / 32))
 cuLaunchKernel(mandel_ker, gridsize, 1, 1, 32, 1, 1, 10000, None, mandel_params, None)
 cuCtxSynchronize()

我们现在将使用cuMemcpyDtoH将数据从 GPU 复制到我们的 NumPy 数组中，使用 NumPy 的array.ctypes.data成员，这是一个 C 指针，将允许我们直接从 C 中访问数组作为堆内存的一部分。我们将使用 Ctypes 的类型转换函数cast将其转换为c_void_p：

 cuMemcpyDtoH( cast(graph.ctypes.data, c_void_p), graph_gpu,  c_size_t(lattice.size**2 *sizeof(c_float)))

我们现在完成了！让我们释放在 GPU 上分配的数组，并通过销毁当前上下文来结束我们的 GPU 会话。然后我们将把图形 NumPy 数组返回给调用函数：

 cuMemFree(lattice_gpu)
 cuMemFree(graph_gpu)
 cuCtxDestroy(cuContext)
 return graph

现在我们可以像以前一样设置我们的main函数：

if __name__ == '__main__':
 t1 = time()
 mandel = mandelbrot(512,-2,2,256, 2)
 t2 = time()
 mandel_time = t2 - t1
 print 'It took %s seconds to calculate the Mandelbrot graph.' % mandel_time

 fig = plt.figure(1)
 plt.imshow(mandel, extent=(-2, 2, -2, 2))
 plt.show()

现在尝试运行这个函数，确保它产生与我们刚刚编写的其他曼德布洛特程序相同的输出。

恭喜你——你刚刚编写了一个直接接口到低级 CUDA Driver API，并成功使用它启动了一个内核！

这个程序也可以在本书的 GitHub 存储库中的目录下的mandelbrot_driver.py文件中找到。

总结

我们从简要概述 Python Ctypes 库开始了本章，该库用于直接与编译的二进制代码进行接口，特别是用 C/C++编写的动态库。然后，我们看了如何使用 CUDA-C 编写一个启动 CUDA 内核的基于 C 的包装器，然后使用这个包装器间接地从 Python 启动我们的 CUDA 内核，方法是使用 Ctypes 编写一个对这个函数的接口。然后，我们学习了如何将 CUDA 内核编译成 PTX 模块二进制文件，可以将其视为一个带有 CUDA 内核函数的 DLL，并看到如何使用 PyCUDA 加载 PTX 文件并启动预编译的内核。最后，我们编写了一系列 CUDA Driver API 的 Ctypes 包装器，并看到我们如何使用这些包装器执行基本的 GPU 操作，包括从 PTX 文件启动预编译的内核到 GPU 上。

我们现在将进入本书中可能是最技术性的一章：第十一章，《CUDA 性能优化》。在本章中，我们将学习关于 NVIDIA GPU 的一些技术细节，这将帮助我们提高应用程序的性能水平。

问题

1. 假设您使用nvcc将包含主机和内核代码的单个.cu文件编译成 EXE 文件，还编译成 PTX 文件。哪个文件将包含主机函数，哪个文件将包含 GPU 代码？

2. 如果我们使用 CUDA Driver API，为什么要销毁上下文？

3. 在本章开始时，当我们首次看到如何使用 Ctypes 时，请注意在调用printf之前，我们必须将浮点值 3.14 强制转换为 Ctypes 的c_double对象。然而，在本章中我们可以看到许多不需要将类型转换为 Ctypes 的工作案例。你认为为什么printf在这里是一个例外呢？

4. 假设您想要向我们的 Python CUDA Driver 接口模块添加功能以支持 CUDA 流。您将如何在 Ctypes 中表示单个流对象？

5. 为什么在mandelbrot.cu中的函数要使用extern "C"？

6. 再次查看mandelbrot_driver.py。为什么我们在 GPU 内存分配和主机/GPU 内存传输之后不使用cuCtxSynchronize函数，而只在单个内核调用之后使用？

第十一章：CUDA 中的性能优化

在这个倒数第二章中，我们将介绍一些相当高级的 CUDA 功能，可以用于低级性能优化。我们将首先学习动态并行性，它允许内核在 GPU 上启动和管理其他内核，并看看我们如何使用它直接在 GPU 上实现快速排序。我们将学习关于矢量化内存访问，可以用于增加从 GPU 全局内存读取时的内存访问加速。然后我们将看看如何使用 CUDA 原子操作，这些是线程安全的函数，可以在没有线程同步或mutex锁的情况下操作共享数据。我们将学习关于 Warp，它是 32 个或更少线程的基本块，在这些线程可以直接读取或写入彼此的变量，然后简要地涉足 PTX 汇编的世界。我们将通过直接在我们的 CUDA-C 代码中内联编写一些基本的 PTX 汇编来做到这一点，这本身将内联在我们的 Python 代码中！最后，我们将把所有这些小的低级调整结合到一个最终的例子中，我们将应用它们来制作一个快速的求和内核，并将其与 PyCUDA 的求和进行比较。

本章的学习成果如下：

• CUDA 中的动态并行性

• 在 GPU 上使用动态并行性实现快速排序

• 使用矢量化类型加速设备内存访问

• 使用线程安全的 CUDA 原子操作

• 基本的 PTX 汇编

• 将所有这些概念应用于编写性能优化的求和内核

动态并行性

首先，我们将看一下动态并行性，这是 CUDA 中的一个功能，允许内核在 GPU 上启动和管理其他内核，而无需主机的任何交互或输入。这也使得许多通常在 GPU 上可用的主机端 CUDA-C 功能也可用于 GPU，例如设备内存分配/释放，设备到设备的内存复制，上下文范围的同步和流。

让我们从一个非常简单的例子开始。我们将创建一个小内核，覆盖N个线程，每个线程将向终端打印一条简短的消息，然后递归地启动另一个覆盖N-1个线程的内核。这个过程将持续，直到N达到 1。(当然，除了说明动态并行性如何工作之外，这个例子将毫无意义。)

让我们从 Python 中的import语句开始：

from __future__ import division
import numpy as np
from pycuda.compiler import DynamicSourceModule
import pycuda.autoinit

请注意，我们必须导入DynamicSourceModule而不是通常的SourceModule！这是因为动态并行性功能需要编译器设置特定的配置细节。否则，这将看起来和行为像通常的SourceModule操作。现在我们可以继续编写内核：

DynamicParallelismCode='''
__global__ void dynamic_hello_ker(int depth)
{
 printf("Hello from thread %d, recursion depth %d!\\n", threadIdx.x, depth);
 if (threadIdx.x == 0 && blockIdx.x == 0 && blockDim.x > 1)
 {
  printf("Launching a new kernel from depth %d .\\n", depth);
  printf("-----------------------------------------\\n");
  dynamic_hello_ker<<< 1, blockDim.x - 1 >>>(depth + 1);
 }
}'''

在这里需要注意的最重要的一点是：我们必须小心，只有一个线程启动下一个迭代的内核，使用一个良好放置的if语句来检查threadIdx和blockIdx的值。如果我们不这样做，那么每个线程将在每个深度迭代中启动远多于必要的内核实例。另外，注意我们可以以正常方式启动内核，使用通常的 CUDA-C 三重括号表示法——我们不必使用任何晦涩或低级命令来利用动态并行性。

在使用 CUDA 动态并行性功能时，一定要小心避免不必要的内核启动。这可以通过指定的线程启动下一个内核迭代来实现。

现在让我们结束这一切：

dp_mod = DynamicSourceModule(DynamicParallelismCode)
hello_ker = dp_mod.get_function('dynamic_hello_ker')
hello_ker(np.int32(0), grid=(1,1,1), block=(4,1,1))

现在我们可以运行前面的代码，这将给我们以下输出：

这个例子也可以在本书的 GitHub 存储库中的dynamic_hello.py文件中找到。

具有动态并行性的快速排序

现在让我们来看一个稍微有趣和实用的动态并行应用——快速排序算法。实际上，这是一个非常适合并行化的算法，我们将会看到。

让我们先简要回顾一下。快速排序是一种递归和原地排序算法，平均和最佳情况下的性能为O(N log N)，最坏情况下的性能为O(N²)。快速排序通过在未排序的数组中选择一个称为枢轴的任意点，然后将数组分成一个左数组（其中包含所有小于枢轴的点）、一个右数组（其中包含所有等于或大于枢轴的点），枢轴位于两个数组之间。如果一个或两个数组的长度大于 1，那么我们将在一个或两个子数组上再次递归调用快速排序，此时枢轴点已经处于最终位置。

在纯 Python 中，可以使用函数式编程一行代码实现快速排序：

qsort = lambda xs : [] if xs == [] else qsort(filter(lambda x: x < xs[-1] , xs[0:-1])) + [xs[-1]] + qsort(filter(lambda x: x >= xs[-1] , xs[0:-1]))

我们可以看到并行性将在快速排序递归调用右数组和左数组时发挥作用——我们可以看到这将从一个线程在初始大数组上操作开始，但当数组变得非常小时，应该有许多线程在它们上工作。在这里，我们实际上将通过在每个单个线程上启动所有内核来实现这一点！

让我们开始吧，并从导入语句开始。（我们将确保从标准随机模块中导入shuffle函数，以备后面的示例。）

from __future__ import division
import numpy as np
from pycuda.compiler import DynamicSourceModule
import pycuda.autoinit
from pycuda import gpuarray
from random import shuffle

现在我们将编写我们的快速排序内核。我们将为分区步骤编写一个device函数，它将接受一个整数指针、要分区的子数组的最低点和子数组的最高点。此函数还将使用此子数组的最高点作为枢轴。最终，在此函数完成后，它将返回枢轴的最终位置。

DynamicQuicksortCode='''
__device__ int partition(int * a, int lo, int hi)
{
 int i = lo;
 int pivot = a[hi];
 int temp;

 for (int k=lo; k<hi; k++)
 {
  if (a[k] < pivot)
  {
   temp = a[k];
   a[k] = a[i];
   a[i] = temp;
   i++;
  }
 }

 a[hi] = a[i];
 a[i] = pivot;

 return i;
}

现在我们可以编写实现此分区函数的内核，将其转换为并行快速排序。我们将使用 CUDA-C 约定来处理流，这是我们到目前为止还没有见过的：要在 CUDA-C 中的流s中启动内核k，我们使用k<<<grid, block, sharedMemBytesPerBlock, s>>>(...)。通过在这里使用两个流，我们可以确保它们是并行启动的。（考虑到我们不会使用共享内存，我们将把第三个启动参数设置为“0”。）流对象的创建和销毁应该是不言自明的：

__global__ void quicksort_ker(int *a, int lo, int hi)
{

 cudaStream_t s_left, s_right; 
 cudaStreamCreateWithFlags(&s_left, cudaStreamNonBlocking);
 cudaStreamCreateWithFlags(&s_right, cudaStreamNonBlocking);

 int mid = partition(a, lo, hi);

 if(mid - 1 - lo > 0)
   quicksort_ker<<< 1, 1, 0, s_left >>>(a, lo, mid - 1);
 if(hi - (mid + 1) > 0)
   quicksort_ker<<< 1, 1, 0, s_right >>>(a, mid + 1, hi);

 cudaStreamDestroy(s_left);
 cudaStreamDestroy(s_right);

}
'''

现在让我们随机洗牌一个包含 100 个整数的列表，并让我们的内核为我们进行排序。请注意我们如何在单个线程上启动内核：

qsort_mod = DynamicSourceModule(DynamicQuicksortCode)

qsort_ker = qsort_mod.get_function('quicksort_ker')

if __name__ == '__main__':
    a = range(100)
    shuffle(a)

    a = np.int32(a)

    d_a = gpuarray.to_gpu(a)

    print 'Unsorted array: %s' % a

    qsort_ker(d_a, np.int32(0), np.int32(a.size - 1), grid=(1,1,1), block=(1,1,1))

    a_sorted = list(d_a.get())

    print 'Sorted array: %s' % a_sorted

此程序也可以在本书的 GitHub 存储库中的dynamic_quicksort.py文件中找到。

矢量化数据类型和内存访问

现在我们将看一下 CUDA 的矢量化数据类型。这些是标准数据类型的矢量化版本，例如 int 或 double，它们可以存储多个值。32 位类型的矢量化版本的大小最多为 4（例如int2，int3，int4和float4），而 64 位变量只能矢量化为原始大小的两倍（例如double2和long2）。对于大小为 4 的矢量化变量，我们使用 C 的“struct”表示法访问每个单独的元素，成员为x，y，z和w，而对于 3 个成员的变量，我们使用x，y和z，对于 2 个成员的变量，我们只使用x和y。

现在这些可能看起来毫无意义，但这些数据类型可以用来提高从全局内存加载数组的性能。现在，让我们进行一个小测试，看看我们如何可以从整数数组中加载一些 int4 变量，以及从双精度数组中加载 double2 变量——我们将不得不使用 CUDA 的reinterpret_cast运算符来实现这一点：

from __future__ import division
import numpy as np
from pycuda.compiler import SourceModule
import pycuda.autoinit
from pycuda import gpuarray

VecCode='''
__global__ void vec_ker(int *ints, double *doubles) { 

 int4 f1, f2;

 f1 = *reinterpret_cast<int4*>(ints);
 f2 = *reinterpret_cast<int4*>(&ints[4]);

 printf("First int4: %d, %d, %d, %d\\n", f1.x, f1.y, f1.z, f1.w);
 printf("Second int4: %d, %d, %d, %d\\n", f2.x, f2.y, f2.z, f2.w);

 double2 d1, d2;

 d1 = *reinterpret_cast<double2*>(doubles);
 d2 = *reinterpret_cast<double2*>(&doubles[2]);

 printf("First double2: %f, %f\\n", d1.x, d1.y);
 printf("Second double2: %f, %f\\n", d2.x, d2.y);

}'''

请注意，我们必须使用dereference运算符*来设置矢量化变量，以及我们必须通过引用（&ints[4]，&doubles[2]）跳转到下一个地址来加载第二个int4和double2，使用数组上的引用运算符&：

这个例子也可以在本书的 GitHub 存储库中的vectorized_memory.py文件中找到。

线程安全的原子操作

我们现在将学习 CUDA 中的原子操作。原子操作是非常简单的、线程安全的操作，输出到单个全局数组元素或共享内存变量，否则可能会导致竞争条件。

让我们想一个例子。假设我们有一个内核，并且在某个时候我们设置了一个名为x的局部变量跨所有线程。然后我们想找到所有x中的最大值，然后将这个值设置为我们用__shared__ int x_largest声明的共享变量。我们可以通过在每个线程上调用atomicMax(&x_largest, x)来实现这一点。

让我们看一个原子操作的简单例子。我们将为两个实验编写一个小程序：

• 将变量设置为 0，然后为每个线程添加 1

• 找到所有线程中的最大线程 ID 值

让我们首先将tid整数设置为全局线程 ID，然后将全局add_out变量设置为 0。过去，我们会通过一个单独的线程使用if语句来改变变量，但现在我们可以使用atomicExch(add_out, 0)来跨所有线程进行操作。让我们导入并编写到这一点的内核：

from __future__ import division
import numpy as np
from pycuda.compiler import SourceModule
import pycuda.autoinit
from pycuda import gpuarray
import pycuda.driver as drv

AtomicCode='''
__global__ void atomic_ker(int *add_out, int *max_out) 
{

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

 atomicExch(add_out, 0);

应该注意的是，虽然原子操作确实是线程安全的，但它们绝不保证所有线程将同时访问它们，它们可能会在不同的时间由不同的线程执行。这可能会有问题，因为我们将在下一步中修改add_out。这可能会导致add_out在一些线程部分修改后被重置。让我们进行块同步以防止这种情况发生：

 __syncthreads();

现在我们可以使用atomicAdd来为每个线程添加1到add_out，这将给我们总线程数：

 atomicAdd(add_out, 1);

现在让我们通过使用atomicMax来检查tid的最大值是多少。然后我们可以关闭我们的 CUDA 内核：

 atomicMax(max_out, tid);

}
'''

现在我们将添加测试代码；让我们尝试在 1 个包含 100 个线程的块上启动这个。我们这里只需要两个变量，所以我们将不得不分配一些大小为 1 的gpuarray对象。然后我们将打印输出：

atomic_mod = SourceModule(AtomicCode)
atomic_ker = atomic_mod.get_function('atomic_ker')

add_out = gpuarray.empty((1,), dtype=np.int32)
max_out = gpuarray.empty((1,), dtype=np.int32)

atomic_ker(add_out, max_out, grid=(1,1,1), block=(100,1,1))

print 'Atomic operations test:'
print 'add_out: %s' % add_out.get()[0]
print 'max_out: %s' % max_out.get()[0]

现在我们准备运行这个：

这个例子也可以在本书的 GitHub 存储库中的atomic.py文件中找到。

Warp shuffling

我们现在将看看所谓的warp shuffling。这是 CUDA 中的一个特性，允许存在于同一个 CUDA Warp 中的线程通过直接读写对方的寄存器（即它们的本地堆栈空间变量）进行通信，而无需使用shared变量或全局设备内存。Warp shuffling 实际上比其他两个选项快得多，更容易使用。这几乎听起来太好了，所以肯定有一个陷阱——确实，陷阱是这只在存在于同一个 CUDA Warp 上的线程之间起作用，这限制了对大小为 32 或更小的线程组的洗牌操作。另一个限制是我们只能使用 32 位或更小的数据类型。这意味着我们不能在 Warp 中洗牌 64 位长长整数或双精度浮点值。

只有 32 位（或更小）的数据类型可以与 CUDA Warp shuffling 一起使用！这意味着虽然我们可以使用整数、浮点数和字符，但不能使用双精度或长长整数！

在我们继续编码之前，让我们简要回顾一下 CUDA Warps。（在继续之前，您可能希望回顾第六章中名为“Warp Lockstep Property”的部分，“调试和分析您的 CUDA 代码”）。CUDA Warp 是 CUDA 中的最小执行单元，由 32 个线程或更少组成，运行在精确的 32 个 GPU 核心上。就像网格由块组成一样，块同样由一个或多个 Warps 组成，取决于块使用的线程数 - 如果一个块由 32 个线程组成，那么它将使用一个 Warp，如果它使用 96 个线程，它将由三个 Warps 组成。即使 Warp 的大小小于 32，它也被视为完整的 Warp：这意味着只有一个单个线程的块将使用 32 个核心。这也意味着 33 个线程的块将由两个 Warps 和 31 个核心组成。

要记住我们在第六章中看到的内容，“调试和分析您的 CUDA 代码”，Warp 具有所谓的Lockstep Property。这意味着 Warp 中的每个线程将完全并行地迭代每条指令，与 Warp 中的每个其他线程完全一致。也就是说，单个 Warp 中的每个线程将同时执行相同的指令，忽略任何不适用于特定线程的指令 - 这就是为什么要尽量避免单个 Warp 中线程之间的任何分歧。NVIDIA 将这种执行模型称为Single Instruction Multiple Thread，或SIMT。到目前为止，您应该明白为什么我们一直在文本中始终使用 32 个线程的块！

在我们开始之前，我们需要学习另一个术语 - Warp 中的lane是 Warp 内特定线程的唯一标识符，它将介于 0 和 31 之间。有时，这也被称为Lane ID。

让我们从一个简单的例子开始：我们将使用__shfl_xor命令在我们的 Warp 内的所有偶数和奇数编号的 Lanes（线程）之间交换特定变量的值。这实际上非常快速和容易做到，所以让我们编写我们的内核并查看一下：

from __future__ import division
import numpy as np
from pycuda.compiler import SourceModule
import pycuda.autoinit
from pycuda import gpuarray

ShflCode='''
__global__ void shfl_xor_ker(int *input, int * output) {

int temp = input[threadIdx.x];

temp = __shfl_xor (temp, 1, blockDim.x);

output[threadIdx.x] = temp;

}'''

除了__shfl_xor之外，这里的一切对我们来说都很熟悉。这是一个 CUDA 线程如何看待这个函数的方式：这个函数从当前线程接收temp的值作为输入。它对当前线程的二进制 Lane ID 执行一个XOR操作，这将是它的左邻居（如果该线程的 Lane 的最低有效位是二进制中的1）或右邻居（如果最低有效位是二进制中的“0”）。然后将当前线程的temp值发送到其邻居，同时检索邻居的 temp 值，这就是__shfl_xor。这将作为输出返回到temp中。然后我们在输出数组中设置值，这将交换我们的输入数组值。

现在让我们编写剩下的测试代码，然后检查输出：

shfl_mod = SourceModule(ShflCode)
shfl_ker = shfl_mod.get_function('shfl_xor_ker')

dinput = gpuarray.to_gpu(np.int32(range(32)))
doutout = gpuarray.empty_like(dinput)

shfl_ker(dinput, doutout, grid=(1,1,1), block=(32,1,1))

print 'input array: %s' % dinput.get()
print 'array after __shfl_xor: %s' % doutout.get()

上述代码的输出如下：

在我们继续之前，让我们做一个更多的 Warp-shuffling 示例 - 我们将实现一个操作，对 Warp 中所有线程的单个本地变量进行求和。让我们回顾一下第四章中的 Naive Parallel Sum 算法，“内核、线程、块和网格”，这个算法非常快速，但做出了一个天真的假设，即我们有与数据片段一样多的处理器 - 这是生活中为数不多的几种情况之一，我们实际上会有这么多处理器，假设我们正在处理大小为 32 或更小的数组。我们将使用__shfl_down函数在单个 Warp 中实现这一点。__shfl_down接受第一个参数中的线程变量，并通过第二个参数中指示的步数移动变量在线程之间，而第三个参数将指示 Warp 的总大小。

让我们立即实现这个。再次，如果您不熟悉 Naive Parallel Sum 或不记得为什么这应该起作用，请查看第四章，内核、线程、块和网格。我们将使用__shfl_down实现一个直接求和，然后在包括整数 0 到 31 的数组上运行这个求和。然后我们将与 NumPy 自己的sum函数进行比较，以确保正确性：

from __future__ import division
import numpy as np
from pycuda.compiler import SourceModule
import pycuda.autoinit
from pycuda import gpuarray

ShflSumCode='''
__global__ void shfl_sum_ker(int *input, int *out) {

 int temp = input[threadIdx.x];

 for (int i=1; i < 32; i *= 2)
     temp += __shfl_down (temp, i, 32);

 if (threadIdx.x == 0)
     *out = temp;

}'''

shfl_mod = SourceModule(ShflSumCode)
shfl_sum_ker = shfl_mod.get_function('shfl_sum_ker')

array_in = gpuarray.to_gpu(np.int32(range(32)))
out = gpuarray.empty((1,), dtype=np.int32)

shfl_sum_ker(array_in, out, grid=(1,1,1), block=(32,1,1))

print 'Input array: %s' % array_in.get()
print 'Summed value: %s' % out.get()[0]
print 'Does this match with Python''s sum? : %s' % (out.get()[0] == sum(array_in.get()) )

这将给我们以下输出：

本节中的示例也可在本书 GitHub 存储库的Chapter11目录下的shfl_sum.py和shfl_xor.py文件中找到。

内联 PTX 汇编

我们现在将初步了解编写 PTX（Parallel Thread eXecution）汇编语言，这是一种伪汇编语言，适用于所有 Nvidia GPU，反过来由即时（JIT）编译器编译为特定 GPU 的实际机器代码。虽然这显然不是用于日常使用，但如果必要，它将让我们在比 C 甚至更低的级别上工作。一个特定的用例是，如果没有其他源代码可用，您可以轻松地反汇编 CUDA 二进制文件（主机端可执行文件/库或 CUDA .cubin 二进制文件）并检查其 PTX 代码。这可以在 Windows 和 Linux 中使用cuobjdump.exe -ptx cuda_binary命令来完成。

如前所述，我们将只涵盖从 CUDA-C 内部使用 PTX 的一些基本用法，它具有特定的语法和用法，类似于在 GCC 中使用内联主机端汇编语言。让我们开始编写我们的 GPU 代码：

from __future__ import division
import numpy as np
from pycuda.compiler import SourceModule
import pycuda.autoinit
from pycuda import gpuarray

PtxCode='''

我们将通过将代码编写到单独的设备函数中进行几个小实验。让我们从一个简单的函数开始，将一个输入变量设置为零。（我们可以在 CUDA 中使用 C++的传址运算符&，我们将在device函数中使用它。）

__device__ void set_to_zero(int &x)
{
 asm("mov.s32 %0, 0;" : "=r"(x));
}

让我们在继续之前先分解一下。asm当然会告诉nvcc编译器我们将使用汇编，所以我们必须将代码放入引号中，以便正确处理。mov指令只是复制一个常量或其他值，并将其输入到寄存器中。（寄存器是 GPU 或 CPU 用于存储或操作值的最基本类型的芯片存储单元；这是 CUDA 中大多数本地变量的使用方式。）mov.s32中的.s32部分表示我们正在使用带符号的 32 位整数变量——PTX 汇编没有 C 中数据的类型，因此我们必须小心使用正确的特定操作。%0告诉nvcc使用与此处字符串的第0个参数对应的寄存器，并用逗号将此与mov的下一个输入分隔开，这是常量0。然后我们以分号结束汇编行，就像我们在 C 中一样，并用引号关闭这个汇编代码字符串。然后我们将使用冒号（而不是逗号！）来指示我们想要在我们的代码中使用的变量。"=r"表示两件事：=将告诉nvcc寄存器将被写入为输出，而r表示这应该被处理为 32 位整数数据类型。然后我们将要由汇编器处理的变量放在括号中，然后关闭asm，就像我们对任何 C 函数一样。

所有这些都是为了将一个变量的值设置为 0！现在，让我们编写一个小的设备函数，用于为我们添加两个浮点数：

__device__ void add_floats(float &out, float in1, float in2)
{
 asm("add.f32 %0, %1, %2 ;" : "=f"(out) : "f"(in1) , "f"(in2));
}

让我们停下来注意一些事情。首先，当然，我们使用add.f32来表示我们要将两个 32 位浮点值相加。我们还使用"=f"表示我们将写入一个寄存器，f表示我们将只从中读取。还要注意我们如何使用冒号来分隔write寄存器和only read寄存器，以供nvcc使用。

在继续之前，让我们再看一个简单的例子，即类似于 C 中的++运算符的函数，它将整数增加1：

__device__ void plusplus(int &x)
{
 asm("add.s32 %0, %0, 1;" : "+r"(x));
}

首先，请注意我们将“0th”参数用作输出和第一个输入。接下来，请注意我们使用的是+r而不是=r——+告诉nvcc这个寄存器在这个指令中将被读取和写入。

现在我们不会变得更复杂了，因为即使在汇编语言中编写一个简单的if语句也是相当复杂的。但是，让我们看一些更多的示例，这些示例在使用 CUDA Warps 时会很有用。让我们从一个小函数开始，这个函数将给出当前线程的 lane ID；这是非常有用的，实际上比使用 CUDA-C 更直接，因为 lane ID 实际上存储在一个称为%laneid的特殊寄存器中，我们无法在纯 C 中访问它。（请注意代码中我们使用了两个%符号，这将告诉nvcc直接在%laneid引用的汇编代码中使用%，而不是将其解释为asm命令的参数。）

__device__ int laneid()
{
 int id; 
 asm("mov.u32 %0, %%laneid; " : "=r"(id)); 
 return id;
}

现在让我们编写另外两个函数，这些函数对处理 CUDA Warps 将会很有用。请记住，只能使用洗牌命令在 Warp 之间传递 32 位变量。这意味着要在 Warp 上传递 64 位变量，我们必须将其拆分为两个 32 位变量，分别将这两个变量洗牌到另一个线程，然后将这两个 32 位值重新组合成原始的 64 位变量。对于将 64 位双精度拆分为两个 32 位整数，我们可以使用mov.b64命令，注意我们必须使用d来表示 64 位浮点双精度：

请注意我们在以下代码中使用volatile，这将确保这些命令在编译后按原样执行。我们这样做是因为有时编译器会对内联汇编代码进行自己的优化，但对于这样特别敏感的操作，我们希望按照原样执行。

__device__ void split64(double val, int & lo, int & hi)
{
 asm volatile("mov.b64 {%0, %1}, %2; ":"=r"(lo),"=r"(hi):"d"(val));
}

__device__ void combine64(double &val, int lo, int hi)
{
 asm volatile("mov.b64 %0, {%1, %2}; ":"=d"(val):"r"(lo),"r"(hi));
}

现在让我们编写一个简单的内核，用于测试我们编写的所有 PTX 汇编设备函数。然后我们将它启动在一个单个线程上，以便我们可以检查一切：

__global__ void ptx_test_ker() { 

 int x=123;

 printf("x is %d \\n", x);

 set_to_zero(x);

 printf("x is now %d \\n", x);

 plusplus(x);

 printf("x is now %d \\n", x);

 float f;

 add_floats(f, 1.11, 2.22 );

 printf("f is now %f \\n", f);

 printf("lane ID: %d \\n", laneid() );

 double orig = 3.1415;

 int t1, t2;

 split64(orig, t1, t2);

 double recon;

 combine64(recon, t1, t2);

 printf("Do split64 / combine64 work? : %s \\n", (orig == recon) ? "true" : "false"); 

}'''

ptx_mod = SourceModule(PtxCode)
ptx_test_ker = ptx_mod.get_function('ptx_test_ker')
ptx_test_ker(grid=(1,1,1), block=(1,1,1))

现在我们将运行前面的代码：

此示例也可在本书的 GitHub 存储库中的Chapter11目录下的ptx_assembly.py文件中找到。

性能优化的数组求和

对于本书的最后一个示例，我们将为给定的双精度数组制作一个标准的数组求和内核，但这一次我们将使用本章学到的所有技巧，使其尽可能快速。我们将检查我们求和内核的输出与 NumPy 的sum函数，然后我们将使用标准的 Python timeit函数运行一些测试，以比较我们的函数与 PyCUDA 自己的gpuarray对象的sum函数相比如何。

让我们开始导入所有必要的库，然后从一个laneid函数开始，类似于我们在上一节中使用的函数：

from __future__ import division
import numpy as np
from pycuda.compiler import SourceModule
import pycuda.autoinit
from pycuda import gpuarray
import pycuda.driver as drv
from timeit import timeit

SumCode='''
__device__ void __inline__ laneid(int & id)
{
 asm("mov.u32 %0, %%laneid; " : "=r"(id)); 
}

让我们注意一些事情——注意我们在设备函数的声明中放了一个新的内联语句。这将有效地将我们的函数变成一个宏，当我们从内核中调用这个函数时，会减少一点时间。另外，注意我们通过引用设置了id变量，而不是返回一个值——在这种情况下，实际上可能有两个整数寄存器应该被使用，并且应该有一个额外的复制命令。这样可以确保这种情况不会发生。

让我们以类似的方式编写其他设备函数。我们需要再编写两个设备函数，以便我们可以将 64 位双精度数拆分和合并为两个 32 位变量：

__device__ void __inline__ split64(double val, int & lo, int & hi)
{
 asm volatile("mov.b64 {%0, %1}, %2; ":"=r"(lo),"=r"(hi):"d"(val));
}

__device__ void __inline__ combine64(double &val, int lo, int hi)
{
 asm volatile("mov.b64 %0, {%1, %2}; ":"=d"(val):"r"(lo),"r"(hi));
}

让我们开始编写内核。我们将接收一个名为 input 的双精度数组，然后将整个总和输出到out，out应该初始化为0。我们将首先获取当前线程的 lane ID，并使用矢量化内存加载将两个值从全局内存加载到当前线程中：

__global__ void sum_ker(double *input, double *out) 
{

 int id;
 laneid(id);

 double2 vals = *reinterpret_cast<double2*> ( &input[(blockDim.x*blockIdx.x + threadIdx.x) * 2] );

现在让我们将双精度变量vals中的这些值求和到一个新的双精度变量sum_val中，它将跟踪本线程的所有求和。我们将创建两个 32 位整数s1和s2，我们将使用它们来分割这个值，并使用 Warp Shuffling 与其他线程共享一个temp变量来重构值：

 double sum_val = vals.x + vals.y;

 double temp;

 int s1, s2;

现在让我们再次在 Warp 上使用 Naive Parallel 求和，这将与在 Warp 上对 32 位整数求和相同，只是我们将在每次迭代中使用sum_val和temp上的split64和combine64PTX 函数：

 for (int i=1; i < 32; i *= 2)
 {

     // use PTX assembly to split
     split64(sum_val, s1, s2);

     // shuffle to transfer data
     s1 = __shfl_down (s1, i, 32);
     s2 = __shfl_down (s2, i, 32);

     // PTX assembly to combine
     combine64(temp, s1, s2);
     sum_val += temp;
 }

现在我们完成了，让我们让每个 Warp 的0th线程将其结束值添加到out，使用线程安全的atomicAdd：

 if (id == 0)
     atomicAdd(out, sum_val);

}'''

我们现在将编写我们的测试代码，使用timeit操作来测量我们的内核和 PyCUDA 对 10000232 个双精度数组进行 20 次迭代的平均时间：

sum_mod = SourceModule(SumCode)
sum_ker = sum_mod.get_function('sum_ker')

a = np.float64(np.random.randn(10000*2*32))
a_gpu = gpuarray.to_gpu(a)
out = gpuarray.zeros((1,), dtype=np.float64)

sum_ker(a_gpu, out, grid=(int(np.ceil(a.size/64)),1,1), block=(32,1,1))
drv.Context.synchronize()

print 'Does sum_ker produces the same value as NumPy\'s sum (according allclose)? : %s' % np.allclose(np.sum(a) , out.get()[0])

print 'Performing sum_ker / PyCUDA sum timing tests (20 each)...'

sum_ker_time = timeit('''from __main__ import sum_ker, a_gpu, out, np, drv \nsum_ker(a_gpu, out, grid=(int(np.ceil(a_gpu.size/64)),1,1), block=(32,1,1)) \ndrv.Context.synchronize()''', number=20)
pycuda_sum_time = timeit('''from __main__ import gpuarray, a_gpu, drv \ngpuarray.sum(a_gpu) \ndrv.Context.synchronize()''', number=20)

print 'sum_ker average time duration: %s, PyCUDA\'s gpuarray.sum average time duration: %s' % (sum_ker_time, pycuda_sum_time)
print '(Performance improvement of sum_ker over gpuarray.sum: %s )' % (pycuda_sum_time / sum_ker_time)

让我们从 IPython 中运行这个。确保你已经先运行了gpuarray.sum和sum_ker，以确保我们不会计算nvcc的编译时间：

因此，虽然求和通常相当无聊，但我们可以因为我们巧妙地利用硬件技巧来加速这样一个单调乏味的算法而感到兴奋。

此示例可在本书的 GitHub 存储库的Chapter11目录下的performance_sum_ker.py文件中找到。

总结

我们开始这一章时学习了动态并行性，这是一种允许我们直接在 GPU 上从其他内核启动和管理内核的范式。我们看到了我们如何可以使用这个来直接在 GPU 上实现快速排序算法。然后我们学习了 CUDA 中的矢量化数据类型，并看到了我们如何可以使用这些类型来加速从全局设备内存中读取数据。然后我们学习了 CUDA Warps，这是 GPU 上的小单位，每个 Warp 包含 32 个线程或更少，并且我们看到了单个 Warp 内的线程如何可以直接读取和写入彼此的寄存器，使用 Warp Shuffling。然后我们看了一下如何在 PTX 汇编中编写一些基本操作，包括确定 lane ID 和将 64 位变量分割为两个 32 位变量等导入操作。最后，我们通过编写一个新的性能优化求和内核来结束了这一章，该内核用于双精度数组，应用了本章学到的几乎所有技巧。我们看到，这实际上比双精度数组长度为 500,000 的标准 PyCUDA 求和更快。

我们已经完成了本书的所有技术章节！你应该为自己感到骄傲，因为现在你肯定是一个技艺高超的 GPU 程序员。现在我们将开始最后一章，在这一章中，我们将简要介绍一些不同的路径，可以帮助你应用和扩展你的 GPU 编程知识。

问题

1. 在原子操作示例中，尝试在启动内核之前将网格大小从 1 更改为 2，同时保持总块大小为 100。如果这给出了add_out的错误输出（除了 200 以外的任何值），那么为什么是错误的，考虑到atomicExch是线程安全的呢？

2. 在原子操作示例中，尝试移除__syncthreads，然后在原始参数的网格大小为 1，块大小为 100 的情况下运行内核。如果这给出了add_out的错误输出（除了 100 以外的任何值），那么为什么是错误的，考虑到atomicExch是线程安全的呢？

3. 为什么我们不必使用__syncthreads来同步大小为 32 或更小的块？

4. 我们发现sum_ker对于长度为 640,000（10000*2*32）的随机值数组比 PyCUDA 的求和操作快大约五倍。如果你尝试在这个数字的末尾加上一个零（也就是乘以 10），你会注意到性能下降到sum_ker只比 PyCUDA 的求和快大约 1.5 倍的程度。如果你在这个数字的末尾再加上一个零，你会注意到sum_ker只比 PyCUDA 的求和快 75%。你认为这是为什么？我们如何改进sum_ker以在更大的数组上更快？

5. 哪种算法执行了更多的加法操作（计算 C +运算符的调用和将 atomicSum 视为单个操作）：sum_ker还是 PyCUDA 的sum？

第十二章：从这里往哪里走

这本书就像一次冒险的登山旅程一样……但现在，最终，我们已经到达了我们的徒步旅行的终点。我们现在站在介绍性 GPU 编程的山顶上，我们骄傲地回望我们的故乡串行编程村，微笑着想着我们旧的一维编程传统的天真，我们曾认为在 Unix 中fork一个进程就是我们对并行编程概念的全部理解。我们经历了许多险阻和危险才到达这一点，我们甚至可能犯了一些错误，比如在 Linux 中安装了一个损坏的 NVIDIA 驱动模块，或者在父母家度假时通过缓慢的 100k 连接下载了错误的 Visual Studio 版本。但这些挫折只是暂时的，留下的伤口变成了使我们更加强大对抗（GPU）自然力量的老茧。

然而，在我们的眼角，我们可以看到离我们站立的地方几米远处有两个木制标志；我们把目光从我们过去的小村庄移开，现在看着它们。第一个标志指向我们当前所面对的方向，上面只有一个词——过去。另一个指向相反的方向，也只有一个词——未来。我们转身朝着指向未来的方向走去，看到一个大而闪亮的大都市展现在我们面前，一直延伸到地平线，招手着我们。现在我们终于喘过气来，可以开始走向未来了…

在本章中，我们将介绍一些你现在可以继续学习 GPU 编程相关领域的选项。无论你是想要建立一个职业生涯，作为一个业余爱好者为了乐趣而做这个，作为一个工程学生为了课程而学习 GPU，作为一个程序员或工程师试图增强你的技术背景，还是作为一个学术科学家试图将 GPU 应用到一个研究项目中，你现在都有很多选择。就像我们比喻的大都市一样，迷失其中很容易，很难确定我们应该去哪里。我们希望在这最后一章中提供类似于简短的导游，为你提供一些下一步可以去的选项。

我们现在将在本章中看一下以下路径：

• 高级 CUDA 和 GPGPU 编程

• 图形

• 机器学习和计算机视觉

• 区块链技术

进一步学习 CUDA 和 GPGPU 编程的知识

你首先可以选择的是，当然是学习更多关于 CUDA 和通用 GPU（GPGPU）编程的知识。在这种情况下，你可能已经找到了一个很好的应用，并且想要编写更高级或优化的 CUDA 代码。你可能对它本身感兴趣，或者你想找一份 CUDA/GPU 程序员的工作。有了这本书提供的坚实的 GPU 编程基础，我们现在将看一些这个领域的高级主题，我们现在已经准备好学习了。

多 GPU 系统

首先想到的一个主要话题是学习如何为安装了多个 GPU 的系统编程。许多专业工作站和服务器都安装了多个 GPU，目的是处理需要不止一个，而是几个顶级 GPU 的数据。为此，存在一个称为多 GPU 编程的子领域。其中大部分工作集中在负载平衡上，即使用每个 GPU 的最大容量，确保没有一个 GPU 被过多的工作饱和，而另一个则未被充分利用。另一个话题是 Inter-GPU 通信，通常涉及一个 GPU 直接使用 CUDA 的 GPUDirect 点对点（P2P）内存访问将内存数组复制到另一个 GPU 或从另一个 GPU 复制。

NVIDIA 在这里提供了有关多 GPU 编程的简要介绍：www.nvidia.com/docs/IO/116711/sc11-multi-gpu.pdf。

集群计算和 MPI

另一个主题是集群计算，即编写程序，使其集体利用包含 GPU 的多台服务器。这些是服务器农场，它们在著名互联网公司（如 Facebook 和 Google）的数据处理设施中，以及政府和军方使用的科学超级计算设施中广泛存在。集群通常使用一种称为消息传递接口（MPI）的编程范式，它是与诸如 C++或 Fortran 之类的语言一起使用的接口，允许您编程连接到同一网络的许多计算机。

有关在 MPI 中使用 CUDA 的更多信息，请参阅此处：devblogs.nvidia.com/introduction-cuda-aware-mpi/。

OpenCL 和 PyOpenCL

CUDA 并不是唯一可以用来编程 GPU 的语言。CUDA 最主要的竞争对手是称为开放计算语言（Open Computing Language）或 OpenCL。CUDA 是一个封闭的专有系统，只能在 NVIDIA 硬件上运行，而 OpenCL 是一个由非营利性 Khronos Group 开发和支持的开放标准。OpenCL 不仅可以用于编程 NVIDIA GPU，还可以用于 AMD Radeon GPU 甚至 Intel HD GPU - 大多数主要技术公司都承诺在其产品中支持 OpenCL。此外，PyCUDA 的作者，UIUC 的 Andreas Kloeckner 教授，还编写了另一个出色（且免费）的 Python 库，名为 PyOpenCL，它提供了一个同样用户友好的 OpenCL 接口，几乎与 PyCUDA 具有相同的语法和概念。

有关 OpenCL 的信息由 NVIDIA 在这里提供：developer.nvidia.com/opencl。

Andreas Kloeckner 的网站上提供了有关免费 PyOpenCL 库的信息：

mathema.tician.de/software/pyopencl/。

图形

显然，GPU 中的 G 代表图形，而在本书中我们并没有看到太多关于图形的内容。尽管机器学习应用现在是 NVIDIA 的支柱产业，但一切都始于渲染出漂亮的图形。我们将在这里提供一些资源，让您开始学习，无论您是想开发视频游戏引擎、渲染 CGI 电影，还是开发 CAD 软件。CUDA 实际上可以与图形应用程序并驾齐驱，并且实际上已经用于专业软件，如 Adobe 的 Photoshop 和 After Effects，以及许多最近的视频游戏，如Mafia和Just Cause系列。我们将简要介绍一些您可能考虑从这里开始的主要 API。

OpenGL

OpenGL 是一个行业开放标准，自 90 年代初就存在。虽然在某些方面它显得有些陈旧，但它是一个稳定的 API，得到了广泛支持，如果你编写了一个使用它的程序，几乎可以保证在任何相对现代的 GPU 上都能运行。CUDA 示例文件夹实际上包含了许多示例，说明了 OpenGL 如何与 CUDA 接口（特别是在2_Graphics子目录中），因此感兴趣的读者可以考虑查看这些示例。（在 Windows 中，默认位置为C:\ProgramData\NVIDIA Corporation\CUDA Samples，在 Linux 中为/usr/local/cuda/samples。）

有关 OpenGL 的信息可以直接从 NVIDIA 这里获取：developer.nvidia.com/opengl。

PyCUDA 还提供了与 NVIDIA OpenGL 驱动程序的接口。有关信息，请访问此处：documen.tician.de/pycuda/gl.html。

DirectX 12

DirectX 12 是微软著名且得到良好支持的图形 API 的最新版本。虽然这是 Windows PC 和微软 Xbox 游戏机的专有，但这些系统显然拥有数亿用户的广泛安装基础。此外，除了 NVIDIA 显卡，Windows PC 还支持各种 GPU，并且 Visual Studio IDE 提供了很好的易用性。DirectX 12 实际上支持低级 GPGPU 编程类型的概念，并且可以利用多个 GPU。

微软的 DirectX 12 编程指南可以在这里找到：docs.microsoft.com/en-us/windows/desktop/direct3d12/directx-12-programming-guide。

Vulkan

Vulkan 可以被认为是 DirectX 12 的开放等效物，由 Khronos Group 开发，作为 OpenGL 的next-gen继任者。除了 Windows，Vulkan 也支持 macOS 和 Linux，以及索尼 PlayStation 4，任天堂 Switch 和 Xbox One 游戏机。Vulkan 具有与 DirectX 12 相同的许多功能，例如准 GPGPU 编程。Vulkan 对 DirectX 12 提供了一些严肃的竞争，例如 2016 年的《毁灭战士》重制版。

Khronos Group 在这里提供了《Vulkan 初学者指南》：www.khronos.org/blog/beginners-guide-to-vulkan。

机器学习和计算机视觉

当然，本章的重点是机器学习及其兄弟计算机视觉。不用说，机器学习（特别是深度神经网络和卷积神经网络的子领域）是如今让 NVIDIA 首席执行官黄仁勋有饭吃的东西。（好吧，我们承认这是本十年的轻描淡写……）如果你需要提醒为什么 GPU 在这个领域如此适用和有用，请再看一下第九章，实现深度神经网络。大量的并行计算和数学运算，以及用户友好的数学库，使 NVIDIA GPU 成为机器学习行业的硬件支柱。

基础知识

虽然现在你已经了解了低级 GPU 编程的许多复杂性，但你不会立即将这些知识应用到机器学习中。如果你在这个领域没有基本技能，比如如何对数据集进行基本统计分析，你真的应该停下来熟悉一下。斯坦福大学教授 Andrew Ng，谷歌 Brain 的创始人，在网上和 YouTube 上提供了许多免费的材料。Ng 教授的工作通常被认为是机器学习教育材料的金标准。

Ng 教授在这里提供了一门免费的机器学习入门课程：www.ml-class.org。

cuDNN

NVIDIA 提供了一个针对深度神经网络基元的优化 GPU 库，名为 cuDNN。这些基元包括前向传播、卷积、反向传播、激活函数（如 sigmoid、ReLU 和 tanh）和梯度下降。cuDNN 是大多数主流深度神经网络框架（如 Tensorflow）在 NVIDIA GPU 上的后端使用。这是 NVIDIA 免费提供的，但必须单独从 CUDA Toolkit 下载。

有关 cuDNN 的更多信息可以在这里找到：developer.nvidia.com/cudnn。

Tensorflow 和 Keras

Tensorflow 当然是谷歌著名的神经网络框架。这是一个免费的开源框架，可用于 Python 和 C++，自 2015 年以来一直向公众提供。

Tensorflow 的教程可以在 Google 这里找到：www.tensorflow.org/tutorials/。

Keras 是一个更高级的库，为 Tensorflow 提供了更用户友好的接口，最初由 Google Brain 的 Francois Chollet 编写。读者实际上可以考虑从 Keras 开始，然后再转向 Tensorflow。

有关 Keras 的信息在这里：keras.io/。

Chainer

Chainer 是另一个神经网络 API，由目前在日本东京大学攻读博士学位的 Seiya Tokui 开发。虽然它比 Tensorflow 更不为人知，但由于其令人难以置信的速度和效率而备受尊重。此外，读者可能会对 Chainer 特别感兴趣，因为最初是使用 PyCUDA 开发的。（后来改用了 CuPy，这是一个为了提供更类似于 NumPy 的接口而开发的 PyCUDA 分支。）

有关 Chainer 的信息在这里：chainer.org/。

OpenCV

自 2001 年以来，开源计算机视觉库（OpenCV）一直存在。该库提供了许多来自经典计算机视觉和图像处理的工具，在深度神经网络时代仍然非常有用。近年来，OpenCV 中的大多数算法都已移植到 CUDA，并且与 PyCUDA 接口非常容易。

有关 OpenCV 的信息在这里：opencv.org/。

区块链技术

最后但并非最不重要的是区块链技术。这是支持比特币和以太坊等加密货币的基础加密技术。这当然是一个非常新的领域，最初由比特币神秘创造者中本聪在 2008 年发表的白皮书中描述。在其发明后几乎立即应用了 GPU——生成货币单位归结为蛮力破解加密谜题，而 GPU 可以并行尝试蛮力破解比一般公众今天可用的任何其他硬件更多的组合。这个过程被称为挖矿。

有兴趣了解区块链技术的人建议阅读中本聪关于比特币的原始白皮书，网址在这里：bitcoin.org/bitcoin.pdf。

GUIMiner 是一个开源的基于 CUDA 的比特币矿工，网址在这里：guiminer.org/。

总结

在本章中，我们介绍了一些对于那些有兴趣进一步了解 GPU 编程背景的选项和路径，这超出了本书的范围。我们介绍的第一条路径是扩展您在纯 CUDA 和 GPGPU 编程方面的背景——您可以学习一些本书未涵盖的内容，包括使用多个 GPU 和网络集群的编程系统。我们还讨论了除 CUDA 之外的一些并行编程语言/API，如 MPI 和 OpenCL。接下来，我们讨论了一些对于有兴趣将 GPU 应用于渲染图形的人来说非常知名的 API，如 Vulkan 和 DirectX 12。然后，我们研究了机器学习，并深入了解了一些您应该具备的基本背景以及一些用于开发深度神经网络的主要框架。最后，我们简要介绍了区块链技术和基于 GPU 的加密货币挖矿。

作为作者，我想对每个人都说谢谢，感谢你们阅读到这里，到了结尾。GPU 编程是我遇到的最棘手的编程子领域之一，我希望我的文字能帮助你掌握基本要点。作为读者，你现在应该可以放纵自己，享用一块你能找到的最丰富、最高热量的巧克力蛋糕——只要知道你赚了它。（但只能吃一块！）

问题

1. 使用谷歌或其他搜索引擎找到至少一个未在本章中介绍的 GPU 编程应用。

2. 尝试找到至少一种编程语言或 API，可以用来编程 GPU，而这在本章中没有提到。

3. 查找谷歌的新张量处理单元（TPU）芯片。这些与 GPU 有何不同？

4. 你认为使用 Wi-Fi 还是有线以太网电缆将计算机连接成集群是更好的主意？

第十三章：评估

第一章，为什么要进行 GPU 编程？

1. 前两个for循环遍历每个像素，它们的输出彼此不变；因此我们可以在这两个for循环上并行化。第三个for循环计算特定像素的最终值，这是固有的递归。

2. 阿姆达尔定律没有考虑在 GPU 和主机之间传输内存所需的时间。

3. 512 x 512 等于 262,144 像素。这意味着第一个 GPU 一次只能计算一半像素的输出，而第二个 GPU 可以一次计算所有像素；这意味着第二个 GPU 在这里将比第一个 GPU 快大约两倍。第三个 GPU 有足够的核心来一次计算所有像素，但正如我们在问题 1 中看到的那样，额外的核心在这里对我们没有用。因此，对于这个问题，第二个和第三个 GPU 的速度将是相同的。

4. 将某个代码段通用地指定为与阿姆达尔定律相关的并行化存在一个问题，即假设这段代码的计算时间在处理器数量N非常大时接近 0。正如我们从上一个问题中看到的，情况并非如此。

5. 首先，一致使用时间可能很麻烦，并且可能无法找出程序的瓶颈。其次，分析器可以告诉您从 Python 的角度来看，所有代码的确切计算时间，因此您可以确定某些库函数或操作系统的后台活动是否有问题，而不是您的代码。

第二章，设置 GPU 编程环境

1. 不，CUDA 只支持 Nvidia GPU，不支持 Intel HD 或 AMD Radeon

2. 本书只使用 Python 2.7 示例

3. 设备管理器

4. lspci

5. free

6. .run

第三章，开始使用 PyCUDA

1. 是的。

2. 主机/设备之间的内存传输和编译时间。

3. 你可以，但这将取决于你的 GPU 和 CPU 设置。

4. 使用 C ?运算符来执行点对点和减少操作。

5. 如果gpuarray对象超出范围，其析构函数将被调用，这将自动在 GPU 上释放它所代表的内存。

6. ReductionKernel可能执行多余的操作，这可能取决于底层 GPU 代码的结构。中性元素将确保没有值因这些多余的操作而改变。

7. 我们应该将neutral设置为带符号 32 位整数的最小可能值。

第四章，内核，线程，块和网格

1. 试试看。

2. 并非所有线程都同时在 GPU 上运行。就像 CPU 在操作系统中在任务之间切换一样，GPU 的各个核心在内核的不同线程之间切换。

3. O(n/640 log n)，即 O(n log n)。

4. 试试看。

5. 实际上，CUDA 中没有内部网格级同步，只有块级（使用__syncthreads）。我们必须将单个块以上的任何内容与主机同步。

6. 天真：129 次加法运算。高效工作：62 次加法运算。

7. 同样，如果我们需要在大量块的网格上进行同步，我们不能使用__syncthreads。如果我们在主机上进行同步，我们还可以在每次迭代中启动更少的线程，从而为其他操作释放更多资源。

8. 在天真的并行求和的情况下，我们可能只会处理一小部分数据点，这些数据点应该等于或少于 GPU 核心的总数，这些核心可能适合于块的最大大小（1032）；由于单个块可以在内部同步，我们应该这样做。只有当数据点的数量远远大于 GPU 上可用核心的数量时，我们才应该使用高效的工作算法。

第五章，流，事件，上下文和并发性

1. 两者的性能都会提高；随着线程数量的增加，GPU 在两种情况下都达到了峰值利用率，减少了使用流所获得的收益。

2. 是的，你可以异步启动任意数量的内核，并使用cudaDeviceSynchronize进行同步。

3. 打开你的文本编辑器，试试看！

4. 高标准差意味着 GPU 的使用不均匀，有时会过载 GPU，有时会过度利用它。低标准差意味着所有启动的操作通常都运行顺利。

5. i. 主机通常可以处理的并发线程远远少于 GPU。ii. 每个线程都需要自己的 CUDA 上下文。GPU 可能会因为上下文过多而不堪重负，因为每个上下文都有自己的内存空间，并且必须处理自己加载的可执行代码。

第六章，调试和分析 CUDA 代码

1. 在 CUDA 中，内存分配会自动同步。

2. lockstep属性仅在大小为 32 或更小的单个块中有效。在这里，两个块将在没有任何lockstep的情况下正确分歧。

3. 在这里也会发生同样的事情。这个 64 线程块实际上会被分成两个 32 线程的 warp。

4. Nvprof 可以计时单个内核启动、GPU 利用率和流使用；任何主机端的分析器只会看到 CUDA 主机函数的启动。

5. Printf 通常更容易用于规模较小、内联内核相对较短的项目。如果你编写了一个非常复杂的 CUDA 内核，有数千行代码，那么你可能会想使用 IDE 逐行步进和调试你的内核。

6. 这告诉 CUDA 我们要使用哪个 GPU。

7. cudaDeviceSynchronize将确保相互依赖的内核启动和内存复制确实是同步的，并且它们不会在所有必要的操作完成之前启动。

第七章，使用 Scikit-CUDA 库

1. SBLAH 以 S 开头，因此该函数使用 32 位实数浮点数。ZBLEH 以 Z 开头，这意味着它使用 128 位复数浮点数。

2. 提示：设置trans = cublas._CUBLAS_OP['T']

3. 提示：使用 Scikit-CUDA 包装器进行点积，skcuda.cublas.cublasSdot

4. 提示：建立在上一个问题的答案基础上。

5. 你可以将 cuBLAS 操作放在一个 CUDA 流中，并使用该流的事件对象来精确测量 GPU 上的计算时间。

6. 由于输入看起来对 cuFFT 来说很复杂，它将计算所有值作为 NumPy。

7. 暗边是由于图像周围的零缓冲造成的。这可以通过在边缘上镜像图像来缓解，而不是使用零缓冲。

第八章，CUDA 设备函数库和 Thrust

1. 试试看。(它实际上比你想象的更准确。)

2. 一个应用：高斯分布可以用于向样本添加“白噪声”以增加机器学习中的数据集。

3. 不，因为它们来自不同的种子，如果我们将它们连接在一起，这些列表可能会有很强的相关性。如果我们打算将它们连接在一起，我们应该使用相同种子的子序列。

4. 试试看。

5. 提示：记住矩阵乘法可以被看作是一系列矩阵-向量乘法，而矩阵-向量乘法可以被看作是一系列点积。

6. Operator()用于定义实际函数。

第九章，实现深度神经网络

1. 一个问题可能是我们没有对训练输入进行归一化。另一个可能是训练速率太大。

2. 使用较小的训练速率，一组权重可能会收敛得非常缓慢，或者根本不会收敛。

3. 大的训练速率可能导致一组权重过度拟合特定的批处理值或该训练集。此外，它可能导致数值溢出/下溢，就像第一个问题一样。

4. Sigmoid。

5. Softmax。

6. 更多更新。

第十章，使用已编译的 GPU 代码

1. 只有 EXE 文件将包含主机函数，但 PTX 和 EXE 都将包含 GPU 代码。

2. cuCtxDestory。

3. printf带有任意输入参数。(尝试查找printf原型。)

4. 使用 Ctypes 的c_void_p对象。

5. 这将允许我们使用原始名称从 Ctypes 链接到函数。

6. CUDA 会自动同步设备内存分配和设备/主机之间的内存复制。

第十一章，CUDA 性能优化

1. atomicExch是线程安全的事实并不保证所有线程将同时执行此函数（因为在网格中不同的块可以在不同的时间执行）。

2. 大小为 100 的块将在多个 warp 上执行，除非我们使用__syncthreads，否则块内部将不会同步。因此，atomicExch可能会被多次调用。

3. 由于 warp 默认以锁步方式执行，并且大小为 32 或更小的块将使用单个 warp 执行，因此__syncthreads是不必要的。

4. 我们在 warp 内使用天真的并行求和，但除此之外，我们使用atomicAdd进行了许多求和，就像我们使用串行求和一样。虽然 CUDA 自动并行化了许多这些atomicAdd调用，但我们可以通过实现高效的并行求和来减少所需的atomicAdd调用总数。

5. 肯定是sum_ker。很明显 PyCUDA 的求和没有使用与我们相同的硬件技巧，因为我们在较小的数组上表现更好，但通过扩大尺寸，PyCUDA 版本更好的唯一解释是它执行的加法操作更少。

第十二章，下一步去哪里

1. 两个例子：DNA 分析和物理模拟。

2. 两个例子：OpenACC，Numba。

3. TPU 仅用于机器学习操作，缺少呈现图形所需的组件。

4. 以太网。