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

Python CUDA GPU 编程实用指南（全）

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

译者：飞龙

协议：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 编程？，给出了我们学习这个领域的一些动机，以及如何应用 Amdahl 定律来估计将串行程序转换为利用 GPU 可能带来的性能提升。

第二章，设置您的 GPU 编程环境，解释了如何在 Windows 和 Linux 下设置适合 CUDA 的 Python 和 C++开发环境。

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

第四章，内核、线程、块和网格，教授编写有效的 CUDA 内核的基础知识，这些内核是启动在 GPU 上的并行函数。我们将看到如何编写 CUDA 设备函数（由 CUDA 内核直接调用的“串行”函数），以及了解 CUDA 的抽象网格/块结构及其在启动内核中的作用。

第五章，流、事件、上下文和并发，涵盖了 CUDA 流的观念，这是一个允许我们并发地在 GPU 上启动和同步多个内核的功能。我们将看到如何使用 CUDA 事件来计时内核启动，以及如何创建和使用 CUDA 上下文。

第六章，调试和性能分析您的 CUDA 代码，填补了我们在纯 CUDA C 编程方面的空白，并展示了如何使用 NVIDIA Nsight IDE 进行调试和开发，以及如何使用 NVIDIA 性能分析工具。

第七章，使用 Scikit-CUDA 中的 CUDA 库，通过 Python Scikit-CUDA 模块简要介绍了几个重要的标准 CUDA 库，包括 cuBLAS、cuFFT 和 cuSOLVER。

第八章，CUDA 设备函数库和 Thrust，展示了如何在我们的代码中使用 cuRAND 和 CUDA Math API 库，以及如何使用 CUDA Thrust C++容器。

第九章，深度神经网络实现，作为一个总结，我们学习如何从头开始构建整个深度神经网络，应用我们在文本中学到的许多想法。

第十章，与编译后的 GPU 代码一起工作，展示了如何使用 PyCUDA 和 Ctypes 将我们的 Python 代码与预编译的 GPU 代码接口。

第十一章，CUDA 中的性能优化，教授了一些非常底层的性能优化技巧，特别是与 CUDA 相关，如 warp 混洗、向量内存访问、使用内联 PTX 汇编和原子操作。

第十二章，从这里开始，概述了一些教育职业道路，这些道路将建立在您现在在 GPU 编程方面的坚实基础之上。

为了充分利用这本书

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

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

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

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

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

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

• 三角学（正弦函数：sin、cos、tan 等）

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

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

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

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

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

• 基于 64 位 x86 Intel/AMD 的 PC

• 4 吉字节（GB）或更多的 RAM

• 入门级 NVIDIA GTX 1050 GPU（帕斯卡架构）或更好

读者应该知道，大多数较老的 GPU 将可能适用于本文中的大多数，如果不是所有示例，但本文中的示例仅在 Windows 10 下的 GTX 1050 和 Linux 下的 GTX 1070 上进行了测试。有关设置和配置的详细说明请参阅第二章，设置您的 GPU 编程环境。

下载示例代码文件

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

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

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

2. 选择 SUPPORT 标签。

3. 点击代码下载和勘误表。

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

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

• Windows 上的 WinRAR/7-Zip

• Mac 上的 Zipeg/iZip/UnRarX

• Linux 上的 7-Zip/PeaZip

书籍的代码包也托管在 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

粗体：表示新术语、重要单词或您在屏幕上看到的单词。例如，菜单或对话框中的单词在文本中显示如下。

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

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

联系我们

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

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

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

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

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

评价

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

想了解更多关于 Packt 的信息，请访问packt.com。

第一章：为什么选择 GPU 编程？

结果表明，除了能够为视频游戏渲染图形之外，图形处理单元（GPU）还为普通消费者提供了一个易于访问的大规模并行****计算手段——一个普通人现在可以从当地的电子产品店购买一张价值 2000 美元的现代 GPU 显卡，将其插入家里的 PC，然后几乎立即使用它进行计算能力，这在 5 到 10 年前只有顶级公司和大学的超级计算机实验室才能提供。近年来，GPU 的这种开放可访问性在许多方面都变得明显，这可以通过对新闻的简要观察来揭示——加密货币矿工使用 GPU 生成比特币等数字货币，遗传学家和生物学家使用 GPU 进行 DNA 分析和研究，物理学家和数学家使用 GPU 进行大规模模拟，AI 研究人员现在可以编程 GPU 编写剧本和创作音乐，而像谷歌和 Facebook 这样的主要互联网公司则使用配备 GPU 的服务器农场进行大规模机器学习任务……这个列表可以一直继续下去。

本书的主要目的是让您尽快掌握 GPU 编程，这样您也可以尽快开始使用它们的强大功能，无论您的最终目标是什么。我们的目标是涵盖如何编程 GPU 的核心基本知识，而不是提供 GPU 如何工作的复杂技术细节和图示。在本书的结尾，我们将提供进一步的资源，以便您可以进一步专业化，并应用您对 GPU 的新知识。（关于特定所需的技术知识和硬件的进一步细节将在本节之后提供。）

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

CUDA 的发音是coo-duh，而不是作为缩写C-U-D-A！CUDA 最初代表Compute Unified Device Architecture，但 NVIDIA 已经放弃了这个缩写，现在将 CUDA 作为一个全大写的正确名称。

我们现在将开始我们的 GPU 编程之旅，首先概述Amdahl 定律。Amdahl 定律是一种简单但有效的方法，可以估计通过将程序或算法卸载到 GPU 上所能获得的潜在速度提升；这将帮助我们确定是否值得我们努力重写代码以利用 GPU。然后我们将简要回顾如何使用cProfile模块分析我们的 Python 代码，以帮助我们找到代码中的瓶颈。

本章的学习成果如下：

• 理解 Amdahl 定律

• 在你的代码中应用 Amdahl 定律

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

技术要求

建议为本章安装 Anaconda Python 2.7：

Anaconda 下载

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

使用 Python 和 CUDA 进行 GPU 编程实战

关于先决条件的更多信息，请查看本书的序言；关于软件和硬件要求，请查看使用 Python 和 CUDA 进行 GPU 编程实战的 README 部分。

并行化和 Amdahl 定律

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

GPU 的强大之处在于其核心数量远多于 CPU，这意味着在吞吐量上有了巨大的进步。这里的吞吐量指的是可以同时进行的计算数量。让我们用一个类比来更好地理解这意味着什么。GPU 就像一条非常宽阔的城市道路，设计用来同时处理许多低速行驶的汽车（高吞吐量，高延迟），而 CPU 则像一条狭窄的高速公路，一次只能允许几辆车通过，但可以将每辆单独的汽车更快地送到目的地（低吞吐量，低延迟）。

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

假设你正在建造一栋房子，并且你已经准备好了所有的设计和材料。你雇佣了一个工人，并估计建造这栋房子需要 100 小时。让我们假设这栋特定的房子可以以这种方式建造，即额外雇佣的每个工人都可以完美地分担工作——也就是说，两个工人需要 50 小时，四个工人需要 25 小时，十个工人需要 10 小时来建造这栋房子——建造你房子的所需小时数将是 100 除以你雇佣的工人数量。这是一个可并行化任务的例子。

我们注意到，对于两个工人来说，完成这项任务的速度是单独一个工人独自建造房屋（即串行）的两倍，对于十个工人一起完成（即并行）则是十倍——也就是说，如果 N 是工人的数量，那么它将是 N 倍快。在这种情况下，N 被称为将我们的任务并行化到串行版本中的加速比。

在我们开始编写给定算法的并行版本之前，我们通常首先估算一下并行化将给我们的任务带来的潜在加速比。这有助于我们确定是否值得投入资源和时间来编写程序的并行化版本。因为现实生活比我们给出的例子要复杂得多，所以很明显，我们不可能总是完美地并行化每个程序——大多数时候，只有我们程序的一部分可以很好地并行化，而其余部分则必须串行运行。

使用 Amdahl 定律

我们现在将推导出Amdahl 定律，这是一个简单的算术公式，用于估算将串行程序的一部分代码并行化到多个处理器上可能产生的潜在速度提升。我们将通过继续我们之前的建造房子的类比来完成这个推导。

上次，我们只考虑了房屋的实际物理建造作为整个时间跨度，但现在，我们也将设计房屋的时间纳入建造房屋的时间跨度。假设世界上只有一个人有设计你房屋的能力——那就是你——而你设计你房屋的计划需要 100 小时。没有其他任何人能与你相比的建筑天赋，所以这部分任务根本不可能被其他建筑师分担——也就是说，设计你的房屋将需要 100 小时，无论你拥有多少资源或能雇佣多少人。所以，如果你只有一个劳动力来建造你的房屋，建造整个房屋所需的时间将是 200 小时——100 小时用于设计，100 小时用于单个劳动力建造。如果我们雇佣两个劳动力，这将需要 150 小时——设计房屋的时间将保持为 100 小时，而建造将需要 50 小时。很明显，建造房屋所需的总小时数将是 100 + 100 / N，其中N是我们雇佣的劳动力数量。

现在，让我们退一步思考，如果我们雇佣一个劳动力，建造房屋需要多少时间——我们最终使用这个来确定随着我们雇佣更多劳动力时的加速比；也就是说，这个过程变得有多快。如果我们雇佣一个劳动力，我们会看到设计和建造房屋所需的时间相同——100 小时。所以，我们可以这样说，在设计上花费的时间部分是 0.5（50%），建造房屋所需的时间部分也是 0.5（50%），当然，这两个部分加起来是 1，即 100%。当我们增加劳动力时，我们想要与这个进行比较——如果我们有两个劳动力，建造的时间部分减半，所以与我们的任务原始串行版本相比，这将需要 0.5 + 0.5/2 = 0.75（75%）的时间，0.75 乘以 200 小时是 150 小时，所以我们可以看到这是有效的。此外，我们可以看到，如果我们有N个劳动力，我们可以计算出N个劳动力的并行化建造将需要多少时间，公式是 0.5 + 0.5 / N。

现在，让我们确定通过增加额外劳动力所获得的 加速比。如果我们有两个劳动力，建造房屋需要 75%的时间，我们可以取 0.75 的倒数来确定我们并行化的加速比——也就是说，加速比将是 1 / 0.75，大约是 1.33 倍于只有一个劳动力的情况。在这种情况下，如果我们有N个劳动力，加速比将是 1 / (0.5 + 0.5 / N)。

我们知道，随着我们添加越来越多的劳动者，.5 / N 将非常接近 0，因此我们可以看到，当你并行化这个任务时，你总能得到一个速度提升的上限——即 1 / (.5 + 0) = 2。我们可以用估计的最大速度提升来除以原始串行时间，以确定这个任务将花费的绝对最小时间——200 / 2 = 100 小时。

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

在这种情况下，不可并行化代码的执行时间比例总是1 – p，所以我们只需要知道p。

我们现在可以使用阿姆达尔定律如下计算速度提升：

总结一下，阿姆达尔定律是一个简单的公式，它允许我们大致（非常粗略地）估计一个至少可以部分并行化的程序的可能速度提升。这可以提供一个一般性的想法，即是否值得编写特定串行程序的并行版本，前提是我们知道我们可以并行化多少代码（p），以及我们可以在多少核心上运行我们的并行化代码（N）。

曼德尔布罗特集

我们现在准备看到一个非常标准的并行计算示例，我们将在本文的后面部分再次讨论——一个生成曼德尔布罗特集图像的算法。让我们首先准确地定义我们的意思。

对于给定的复数c，我们定义一个递归序列用于，其中和对于。如果|z[n]|在n增加到无穷大时仍然被 2 所限制，那么我们将说c是曼德尔布罗特集的成员。

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

现在，让我们考虑如何在 Python 中生成这个集合。我们首先必须考虑一些事情——因为我们显然不能检查每个复数是否在 Mandelbrot 集中，我们必须选择一个特定的范围来检查；我们必须确定每个范围内我们将考虑的点数（宽度，高度）；以及我们将检查的最大值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 集并将其输出到文件，并使用 time 函数来计时这两个操作：

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)

现在让我们运行这个程序（这个程序也可以作为mandelbrot0.py文件，位于 GitHub 仓库的1文件夹中）：

生成 Mandelbrot 集大约需要 14.62 秒，而将图像输出大约需要 0.11 秒。正如我们所看到的，我们是逐点生成 Mandelbrot 集的；不同点的值之间没有相互依赖性，因此它是一个本质上可并行化的函数。相比之下，输出图像的代码不能并行化。

现在，让我们从 Amdahl 定律的角度来分析这个问题。如果我们在这里并行化代码，我们能获得什么样的加速效果？整个程序运行大约需要 14.73 秒；由于我们可以并行化 Mandelbrot 集的生成，我们可以认为并行化代码的执行时间部分p = 14.62 / 14.73 = .99。这个程序有 99%的并行性！

我们能获得什么样的加速效果？嗯，我现在正在使用一台配备入门级 GTX 1050 GPU（640 个核心）的笔记本电脑；因此，当我们使用公式时，我们的N将是 640。我们如下计算加速比：

这确实非常好，这会告诉我们，将我们的算法编程以使用 GPU 是值得的。记住，Amdahl 定律只提供了一个非常粗略的估计！当我们将计算卸载到 GPU 时，会有额外的考虑因素，例如 CPU 向 GPU 发送和接收数据所需的时间；或者，卸载到 GPU 的算法只有部分可并行化。

代码性能分析

在前面的例子中，我们看到了我们可以使用 Python 中的标准 time 函数单独计时不同的函数和组件。虽然这种方法对于我们的小型示例程序来说效果不错，但这并不总是适用于调用许多不同函数的大型程序，其中一些可能或可能不值得我们并行化，甚至优化在 CPU 上。我们的目标是找到程序的瓶颈和热点——即使我们感到精力充沛，在每次函数调用周围使用 time，我们可能也会错过某些东西，或者可能有一些系统或库调用我们没有考虑，但恰好是它们减慢了速度。在我们甚至考虑将代码重写为在 GPU 上运行之前，我们应该找到候选的代码部分来卸载到 GPU 上；我们必须始终遵循著名美国计算机科学家唐纳德·克努特（Donald Knuth）的明智话语：过早优化是万恶之源。

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

使用 cProfile 模块

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

这在 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 之间呢？

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 基于的 Ubuntu 版本。

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

Packt 作者 Sebastian Raschka 博士提供了一个 Python 2.7 和 3.x 之间关键差异的列表，请参阅sebastianraschka.com/Articles/2014_python_2_3_key_diff.html。

我们特别建议 Windows 和 Linux 用户使用 Anaconda Python 2.7 发行版，因为它可以基于用户安装，无需sudo或管理员访问权限，包含本文本中所需的所有数据科学和可视化模块，并使用利用 Intel 的数学内核库（MKL）的快速预优化 NumPy/SciPy 包。（默认的 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/）以获取更多详细信息。此外，还有其他硬件的可能性：对嵌入式系统或具有某些板卡经验（如 Raspberry Pi）的机器人感兴趣的读者可能希望从基于 ARM 的 NVIDIA Jetson 开发板开始，而对云计算或网页编程感兴趣的读者可能考虑远程使用适当的 Azure 或 AWS 实例。在这些情况下，鼓励读者阅读官方文档来设置他们的驱动程序、编译器和 CUDA 工具包。本章中的一些步骤可能适用也可能不适用。

本章的学习成果包括：

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

• 安装 NVIDIA GPU 驱动程序

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

• 安装 NVIDIA CUDA 工具包

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

技术要求

对于本章，建议在 www.anaconda.com/download/ 安装 Anaconda Python 2.7。

本章的代码也可在 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 位 Intel/AMD 的 PC

• 4 GB（GB）的 RAM

• NVIDIA GeForce GTX 1050 GPU（或更高版本）

此配置将确保您可以舒适地学习 GPU 编程，运行本书中的所有示例，并且还可以运行一些其他较新且有趣的基于 GPU 的软件，例如 Google 的 TensorFlow（一个机器学习框架）或 Vulkan SDK（一个前沿的图形 API）。

请注意，您必须拥有 NVIDIA 品牌的 GPU 才能使用本书！ CUDA 工具包是 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（帕斯卡）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 键。这将告诉我们第一行第一项中我们拥有的总内存量（以千兆字节为单位），以及下一行中交换空间中的内存量：

这当然足够了。

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

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

检查您的硬件（Windows）

首先，我们必须打开 Windows 控制面板。我们通过按 Windows + R 并在提示符中输入 Control Panel 来完成此操作，如下面的截图所示：

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

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

安装 GPU 驱动程序

如果你已经安装了 GPU 的驱动程序，你可能可以跳过这一步；此外，一些 CUDA 版本已经预包装了最新驱动程序。通常，CUDA 对安装的驱动程序非常挑剔，甚至可能不与 CUDA 工具包驱动程序兼容，因此你可能需要尝试几个不同的驱动程序才能找到可以工作的一个。

通常来说，Windows 的 CUDA 驱动程序兼容性更好，安装过程也更用户友好。Windows 用户可以考虑跳过这一步，直接使用与 CUDA 工具包捆绑的驱动程序，我们将在本章稍后安装。然而，我们强烈建议 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 服务器设置管理器应该会出现，并指示你正在使用适当的驱动程序版本：

安装 GPU 驱动程序（Windows）

再次强调——通常建议读者最初跳过此步骤，然后安装 CUDA 工具包中包含的驱动程序。

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 工具包了。

设置 Visual Studio (Windows)

在撰写本文时，只有 Visual Studio 2015 似乎与 Python 和最新的 CUDA 工具包完美集成；即 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++（如果您需要或需要其他目的的其他软件包或编程语言，请随意选择，但我们只需要 Visual C++ 用于 GPU 编程）：

这可能需要一些时间来安装。完成此步骤后，我们将准备好安装 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 工具包（Windows）

对于 Windows 用户，您可以通过双击 .exe 文件并遵循所有屏幕提示来安装包。

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

设置我们的 Python 环境以进行 GPU 编程

在我们的系统上正确安装了编译器、IDE 和 CUDA 工具包之后，我们现在可以为 GPU 编程设置一个合适的 Python 环境。这里有多种选择，但我们明确建议您使用 Anaconda Python 发行版。Anaconda Python 是一个自包含且用户友好的发行版，可以直接安装到您的用户目录中，并且安装、使用或更新时不需要任何管理员或sudo级别的系统访问权限。

请记住，Anaconda Python 有两种版本——Python 2.7 和 Python 3。由于 Python 3 目前对我们将使用的某些库支持不佳，因此本书中将使用 Python 2.7，它仍然有广泛的主流使用。

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

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

安装 PyCUDA（Linux）

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

PyCUDA 的安装有几种选择。最简单的方法是通过在命令行中输入pip install pycuda从 PyPI 仓库安装最新稳定版本。您也可以通过遵循 PyCUDA 官方网站mathema.tician.de/software/pycuda/上的说明来安装 PyCUDA 的最新版本。请注意，如果您希望从不同的来源重新安装 PyCUDA，请务必先使用pip uninstall pycuda卸载它。

创建环境启动脚本（Windows）

Windows 用户需要特别注意，为了使用 PyCUDA，他们必须确保 Visual Studio 和 Anaconda Python 环境变量设置正确；否则，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 替换为适当的 IDE 启动命令。）

安装 PyCUDA（Windows）

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

pip install pycuda.whl

（或者，您也可以尝试从 PyPI 存储库使用 pip install pycuda 安装 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。

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

恭喜，您现在可以开始进入 GPU 编程的世界了！

摘要

设置 Python 环境以进行 GPU 编程可能是一个非常细致的过程。对于本文本的目的，建议 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 的特性，例如核心数量、架构和内存。然后，我们将花一些时间熟悉如何使用 PyCUDA 的 gpuarray 类在 Python 和 GPU 之间传输内存，以及如何使用此类进行基本计算。本章的其余部分将展示如何编写一些基本函数（我们将它们称为 CUDA 内核），我们可以直接在 GPU 上启动这些函数。

本章的学习成果如下：

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

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

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

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

• 理解 reduce/scan 操作的函数式编程概念以及如何创建基本的 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 目录中包含了一个名为 deviceQuery 的纯 CUDA-C 命令行示例（适用于 Windows 和 Linux），我们可以运行此操作。让我们看看作者在 Windows 10 笔记本电脑（这是一台配备 GTX 1050 GPU 的 Microsoft Surface Book 2）上产生的输出：

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

虽然你可以从deviceQuery中看到许多其他技术细节，但核心数和内存量通常是你在新 GPU 上第一次运行此程序时应该首先关注的前两件事，因为它们可以给你关于新设备容量的最直接的想法。

使用 PyCUDA 查询 GPU

现在，最后，我们将通过编写我们自己的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))

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

    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 设备将具有多个 SMs，每个 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 之间传输数据

正如我们从先前的 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 发送数据时，明确设置数据类型是一个好主意。这样做的原因有两方面：首先，因为我们使用 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 计算时间的比较，看看是否以及何时在 GPU 上执行这些操作有任何优势。

速度测试

让我们编写一个小程序（time_calc0.py），它将在 CPU 上执行标量乘法，然后在 GPU 上执行相同的操作。然后我们将使用 NumPy 的allclose函数来比较两个输出值。我们将生成一个包含 5000 万个随机 32 位浮点数的数组（这将大约是 48 兆字节的数据，所以这应该在任何稍微现代的主机和 GPU 设备上使用几个 GB 的内存都是完全可行的），然后我们将测量在两个设备上对数组进行标量乘以 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 年的 Microsoft Surface Book 2 上运行的，配备 Kaby Lake i7 处理器和 GTX 1050 GPU。）

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

首先，让我们使用以下行在 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 编译器进行编译和链接；然后它被缓存起来，如果代码再次被调用，则不需要重新编译。这甚至可能包括像这样的标量乘法这样的简单操作！（我们最终会看到，这可以通过在第十章，与编译的 GPU 代码一起工作中使用预编译的代码来改善，或者通过使用 Scikit-CUDA 模块与 NVIDIA 自己的线性代数库，我们将在第七章，使用 CUDA 库与 Scikit-CUDA中看到）。

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

使用 PyCUDA 的 ElementWiseKernel 进行逐点计算

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

在本文中，我们经常使用术语内核；通过内核，我们始终指的是由 CUDA 直接在 GPU 上启动的函数。我们将使用 PyCUDA 的几个函数来生成不同类型内核的模板和设计模式，这将有助于我们过渡到 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"），这通常是以 C 指针的形式在 GPU 上分配的内存。在第二行，我们定义了我们的逐元素操作"out[i] = 2*in[i];"，这将把in中的每个点乘以 2，并将其放置在out的相应索引中。

注意 PyCUDA 会自动为我们设置整数索引i。当我们使用i作为索引时，ElementwiseKernel将自动在我们的 GPU 的许多核心之间并行化我们的计算。最后，我们给我们的代码块赋予其内部的 CUDA C 内核名称（"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 在第一次使用nvcc编译器调用给定的 GPU 内核函数时编译我们的内联 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 数组发送到 GPU (host_data)，该函数自动在 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对象的析构函数管理所有内存清理，你不需要在之后释放此内存。

我们现在可以启动内核了；我们唯一需要做的更改是更改

我们还没有编写生成 Mandelbrot 集的内核函数，但让我们先写一下我们希望这个函数的其他部分如何进行：

    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类型），它将指示哪些元素是 Mandelbrot 集的成员，第三个将是一个整数，表示每个点的最大迭代次数，最后一个输入将用于确定每个点是否属于 Mandelbrot 类的上界。请注意，我们对所有输入到 GPU 中的类型转换都非常小心！

下一行将我们从 GPU 生成的 Mandelbrot 集检索回 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 中与数据交互时，曼德布罗特集的二维结构将被保留。）

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

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

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

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

让我们检查导出的图像以确保这是正确的：

这确实是与第一章中产生的相同曼德布罗特图像，因此我们成功将其在 GPU 上实现了！现在让我们看看我们获得的速度提升：在第一章中，我们用了 14.61 秒来生成这个图表；在这里，它只用了 0.894 秒。记住，PyCUDA 还必须在运行时编译和链接我们的 CUDA C 代码，以及将内存从 GPU 传输到和从 GPU 传输回来的时间。尽管如此，即使有所有这些额外的开销，这仍然是一个非常值得的速度提升！（您可以通过 Git 仓库中名为gpu_mandelbrot0.py的文件查看我们 GPU 曼德布罗特的代码。）

函数式编程的简要探索

在我们继续之前，让我们简要回顾一下 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 接受一个列表，对其执行递归二进制操作，并输出一个单例。让我们通过输入 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] 索引检索到它）。

现在，让我们最后看看一个 PyCUDA 函数，用于生成 GPU 内核——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 上索引的。我们按照与 InclusiveScanKernel 相同的方式设置 reduce_expr。这将从逐元素操作的结果中取出输出，并在数组上执行 reduce 类型的操作。最后，我们使用 neutral 设置 中性元素。这是一个在 reduce_expr 中充当恒等元的元素；在这里，我们设置 neutral=0，因为 0 在加法下总是恒等元（在乘法下，1 是恒等元）。我们将在本书后面更深入地介绍并行前缀时看到为什么必须设置这个。

摘要

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

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

问题

1. 在 simple_element_kernel_example0.py 中，我们在测量 GPU 计算的时间时没有考虑从 GPU 到 GPU 的内存传输。尝试使用 Python 的 time 命令测量 gpuarray 函数 to_gpu 和 get 所花费的时间。考虑到内存传输时间，你认为将这个特定函数卸载到 GPU 上是否值得？

2. 在 第一章 为什么进行 GPU 编程？ 中，我们讨论了 Amdahl 的定律，它给我们一些关于通过将程序的部分卸载到 GPU 上可能获得的收益的想法。在本章中，我们看到了两个 Amdahl 的定律没有考虑的问题。

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

4. 使用 ReductionKernel 创建一个内核，该内核接受两个长度相同的 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 的某些底层技术细节（例如核心、战程和流式多处理器，我们将在本书的后续章节中介绍），以及我们如何使用这些概念来减轻并行编程的认知负担。我们将学习关于线程同步（包括块级和网格级）以及 CUDA 中使用 全局 和 共享 内存 的线程间通信。最后，我们将深入了解如何在 GPU 上实现我们自己的并行前缀类型算法的技术细节（即我们在上一章中介绍过的扫描/归约类型函数），这将使我们能够将本章学到的所有原理付诸实践。

本章的学习成果如下：

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

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

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

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

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

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

技术要求

为了本章，需要一个配备现代 NVIDIA GPU（2016 年及以后）的 Linux 或 Windows 10 PC，并安装所有必要的 GPU 驱动程序和 CUDA 工具包（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 内核函数作为内联 CUDA C，并使用 PyCUDA 将它们启动到我们的 GPU 上。在上一章中，我们使用了 PyCUDA 提供的模板来编写符合特定设计模式的内核；相比之下，我们现在将看到如何从头开始编写我们自己的内核，这样我们就可以编写各种灵活的内核，这些内核可能不符合 PyCUDA 覆盖的任何特定设计模式，并且我们可以对内核有更精细的控制。当然，这些收益将伴随着编程复杂性的增加；我们特别需要理解线程、块和网格及其在内核中的作用，以及如何同步执行内核的线程，以及了解如何在线程之间交换数据。

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

PyCUDA 的 SourceModule 函数

我们将使用 PyCUDA 的SourceModule函数来编译原始的内联 CUDA C 代码，将其编译成可用的内核，这样我们就可以从 Python 中启动它们。我们应该注意，SourceModule实际上是将代码编译成一个CUDA 模块，这就像一个 Python 模块或 Windows DLL，只是它包含了一组编译后的 CUDA 代码。这意味着在使用 PyCUDA 的get_function获取我们想要使用的内核的引用之前，我们不得不“提取”出来。让我们从一个基本的例子开始，看看如何使用SourceModule来使用 CUDA 内核。

和之前一样，我们将从一个最简单的内核函数开始——一个将向量乘以标量的函数。我们首先进行导入：

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通过 PyCUDA 设置的一个值i自动并行化多个 GPU 线程；每个单独线程的标识由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。这是因为你可以对三个维度进行索引，而不仅仅是单个维度。我们为什么要这样做呢？让我们回顾一下关于计算曼德布罗特集的例子，来自第一章，为什么进行 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 的含义——这些将是我们在用于 LIFE 的二维格子上单个 CUDA 线程单元的笛卡尔 x 和 y 值。我们将在一个由二维块组成的二维网格上启动内核，这些块将对应整个单元格。我们将必须使用线程和块常量来找到格子上的笛卡尔点。让我们看看一些图表来阐明这一点。一个位于二维 CUDA 块中的线程可以如下可视化：

到目前为止，你可能想知道为什么我们不只在单个块上启动内核，这样我们就可以直接将_X设置为threadIdx.x，将_Y设置为threadIdx.y，然后完成。这是由于 CUDA 对我们施加的块大小限制——目前，仅支持最多由 1,024 个线程组成的块。这意味着我们最多只能制作 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 实现中，我们将端点“环绕”到晶格的另一侧——例如，我们将x-值为-1视为_WIDTH - 1，将y-值为-1视为_HEIGHT - 1，同样，我们将x-值为_WIDTH视为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 设备函数 是一个串行 C 函数，它由内核中的单个 CUDA 线程调用。虽然这些函数本身是串行的，但它们可以通过多个 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模块中的选择函数为我们的晶格设置一个初始状态。我们将使用随机数填充一个N x N的整数图，一和零；一般来说，如果大约 25%的点是一，其余的是零，我们可以生成一些有趣的晶格动画，所以我们就这样做：

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()

我们现在可以运行我们的程序并享受表演（代码也作为conway_gpu.py文件在 GitHub 仓库的4目录下提供）：

线程同步和交互

我们现在将讨论 GPU 编程中的两个重要概念——线程同步和线程交互。有时，在继续任何进一步的计算之前，我们需要确保每个线程都达到了代码中的同一确切行；我们称之为线程同步。同步与线程交互协同工作，即不同的线程相互传递和读取输入；在这种情况下，我们通常想要确保在传递任何数据之前，所有线程都在计算中的同一步骤上对齐。我们将从这里开始，了解 CUDA __syncthreads设备函数，它用于同步内核中的单个块。

使用__syncthreads()设备函数

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

现在让我们假设我们想要做一些稍微不同的事情——我们想要重新编写我们的内核，使其在给定的细胞晶格上执行一定数量的迭代，而不需要主机反复重新启动。这最初可能看起来很微不足道——一个简单的解决方案可能就是添加一个整数参数来指示迭代次数，并在内联conway_ker内核中的for循环中，做一些额外的简单更改，然后完成。

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

我们想在 GPU 上进行多个完全同步的 LIFE 迭代；我们还将希望使用单个内存数组来表示格。我们可以通过使用名为 __syncthreads() 的 CUDA 设备函数来避免竞态条件。这个函数是一个块级同步屏障——这意味着当线程遇到 __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();
 } 
}
""")

我们现在将像之前一样启动内核并显示输出，对晶格进行 1,000,000 次迭代。请注意，由于每个块中线程数量的限制为 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宏。请注意，我们在复制后确保放置__syncthreads()，以确保在继续 LIFE 算法之前，所有对晶格的内存访问都已完全完成：

 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 个和的集合。在串行操作中生成这些 n 个和通常需要 O(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() 来检索它，回忆一下，在 Python 中 "-1" 索引给出数组的最后一个成员。这将 vec 中每个元素的求和；部分和在 outvec_gpu 的先前值中。（这个例子可以在 GitHub 仓库中的 naive_prefix.py 文件中看到。）

从本质上讲，并行前缀算法必须运行在 n 个线程上，对应于大小为 n 的数组，其中 n 是二进制的（再次强调，这意味着 n 是 2 的某个幂）。然而，如果我们假设我们的算子有一个恒等元素（或者说，中性元素），即对于任何 x 值，都有—。在这种情况下，如果我们的算子是 +，则恒等元素是 0；如果它是 ，则恒等元素是 1；我们只是用一系列 e 值填充元素 ，以便我们有一个新的集合的二进制基数 。

包含与排除前缀

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

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

注意，独占算法产生的输出几乎与包含算法相同，只是它向右移动并省略了最终值。因此，只要我们保留 的副本，我们就可以从任一算法中轻易地获得等效的输出。

高效并行前缀算法

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

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

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

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

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

高效并行前缀（上推阶段）

这是上推阶段的伪代码。（注意对 j 变量的 parfor，这意味着此代码块可以并行化到由 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 的任意大小的数组。这意味着这将操作网格以及块；因此，我们将不得不使用主机进行同步；此外，这需要我们实现两个单独的内核，用于 up-sweep 和 down-sweep 阶段，它们将作为两个阶段中的parfor循环，以及 Python 函数，将作为 up-和 down-sweep 的外部for循环。

让我们从 up-sweep 内核开始。由于我们将从主机迭代重新启动此内核，我们还需要一个参数来指示当前迭代（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——这将使我们能够移除伪代码中的"if j is divisible by 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。

我们现在可以用一行代码运行 up-sweep 阶段，注意到j确实可以被 2^(k+1)整除，这是其构造决定的：

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

我们已经完成了内核的编写！但当然，这并不是 up-sweep 的完整实现。我们还需要用 Python 来完成剩余的部分。让我们获取我们的内核并开始实现。这部分主要自说自话，因为它完全遵循伪代码；我们应该记住，我们是通过使用[:]从x_gpu复制来更新x_old_gpu的，这将保留内存分配，并且只是复制新数据而不是重新分配。此外，请注意我们如何根据要启动的线程数量设置我们的块和网格大小——我们试图保持块大小为 32 的倍数（这是我们在本文中的经验法则，我们将在第十一章，CUDA 性能优化中具体说明为什么我们特别使用 32）。我们应该在文件开头放置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()深入研究了块级同步，以及通过使用共享内存进行块级线程间通信；我们还看到单个块可以操作的线程数量有限，因此当我们创建将在更大网格中使用多个块的内核时，我们必须小心使用这些功能。

我们概述了并行前缀算法的理论，并以实现一个简单的并行前缀算法作为结束，该算法作为一个单独的内核可以操作大小受限为 1,024 的数组（它与___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 的数据集，并且我们有 n 个或更多的处理器，原始并行前缀算法的时间复杂度为 O(log 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 内存操作。这意味着我们可以同时启动多个内存和内核操作，而无需等待每个操作完成。然而，另一方面，我们必须要有一定的组织性，以确保所有相互依赖的操作都得到同步；这意味着我们不应该在输入数据完全复制到设备内存之前启动特定的内核，或者不应该在内核执行完成之前将启动内核的输出数据复制到主机。

为了达到这个目的，我们有了所谓的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 仓库中的multi-kernel.py文件下找到，位于5目录中。）

当然，我们首先需要导入适当的 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 上，迭代地在 64 个线程上启动我们的内核，然后将输出数据复制回主机，并使用 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)

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

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

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

事件和流

我们现在将了解如何使用事件对象与流相关联；这将使我们能够对各种 GPU 操作的流程有高度细致的控制，使我们能够确切地知道每个单独的流通过query函数已经进展到什么程度，甚至允许我们在忽略其他流的同时，将特定流与主机同步。

首先，我们必须认识到这一点——每个流都必须有自己的专用事件对象集合；多个流不能共享一个事件对象。让我们通过修改先前的示例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 上下文之间没有确切的 1 对 1 关系）。

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

同步当前上下文

我们将了解如何在 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()

我们现在可以运行这个程序，它将生成 Mandelbrot 图像到磁盘，就像在第三章，“PyCUDA 入门”中一样。

（此示例也可在存储库中作为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()

就这样！我们应该始终记住，在程序退出之前，使用 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 中创建一个可以返回值的单个主机线程的示例。 (此示例也可以在存储库中的 5 目录下的 single_thread_example.py 文件中找到。) 我们将通过使用 threading 模块中的 Thread 类来创建 Thread 的子类来实现这一点，如下所示：

import threading
class PointlessExampleThread(threading.Thread):

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

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

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

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 测试用例。我们将探讨如何编译 CUDA-C 程序，无论是使用命令行中的 nvcc 还是使用 Nsight IDE。然后，我们将了解如何在 Nsight 中进行调试并使用 Nsight 来理解 CUDA 锁步属性。最后，我们将概述 NVIDIA 命令行和 Visual Profilers，以分析我们的代码。

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

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

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

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

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

• 理解 CUDA warp 锁步属性以及为什么我们应该避免单个 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表示"新行"或"返回"，将输出在终端移动到下一行。）如果我们要从 C 中直接打印任何常数或变量，printf也可以接受可变数量的参数：如果我们想打印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."。

现在，这本书已经快读了一半，我们终于要开始创建我们的第一个 CUDA 并行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);

再次运行代码，我们得到以下输出：

有两个问题立即显得突出：行和列元组的值有重复（每个单独的元组应该只表示一次），并且当它们都应该达到三个时，行和列的值从未超过两个。这应该提示我们，我们在计算行和列的值时是错误的；确实，我们忘记将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确实与设备函数声明中的相反。我们修复了这个问题，但当我们再次运行程序时，我们得到了另一个断言错误。

让我们在设备函数中放置另一个printf调用，在for循环内；这当然是要执行矩阵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 可执行文件（我们将在下一节中看到如何使用 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不同），我们必须使用与号操作符，这将是一个指向指针的指针，因为它是一个指向浮点指针的指针，就像这里一样（float **）。由于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 复制到主机，或者使用 cudaMemcpyDeviceToDevice 在 GPU 上的两个数组之间复制。

现在我们将使用另一个 cudaMalloc 调用来分配一个数组，用于在 GPU 上存储矩阵乘法的结果：

cudaMalloc((float **) &d_output, num_bytes);

最后，当我们想要检查内核的输出时，我们将在主机上设置一些内存来存储 GPU 的输出。让我们设置一个常规的 C float 指针，并使用 malloc 分配内存，就像我们通常做的那样：

float * h_output;
h_output = (float *) malloc(num_bytes);

现在，我们几乎准备好启动我们的内核。CUDA 使用一个名为 dim3 的数据结构来指示内核启动的块和网格大小；我们将这样设置它们，因为我们想要一个 2x2 维度的网格和同样维度的块：

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 函数。

现在我们保存我们的文件，并使用命令行中的 nvcc matrix_ker.cu -o matrix_ker 将其编译成 Windows 或 Linux 可执行文件。这应该输出一个二进制可执行文件，matrix_ker.exe（在 Windows 上）或 matrix_ker（在 Linux 上）。让我们现在尝试编译和运行它：

恭喜你，你刚刚创建了你第一个纯 CUDA-C 程序！（此示例作为 matrix_ker.cu 存储在存储库中，位于 7 目录下。）

使用 Nsight IDE 进行 CUDA-C 开发和调试

现在我们来学习如何使用 Nsight IDE 开发 CUDA-C 程序。我们将看到如何导入我们刚刚编写的程序，并在 Nsight 中编译和调试它。请注意，Nsight 的 Windows 和 Linux 版本之间存在差异，因为它在 Windows 下是 Visual Studio IDE 的插件，在 Linux 下是 Eclipse IDE 的插件。我们将在以下两个小节中介绍这两个版本；如果你不适用某个操作系统，可以自由跳过。

在 Windows 中使用 Nsight 与 Visual Studio

打开 Visual Studio，然后点击文件，然后选择新建 | 项目.... 将弹出一个窗口，你可以设置项目的类型：选择 NVIDIA 下拉菜单项，然后选择 CUDA 9.2：

给项目起一个合适的名字，然后点击确定。解决方案资源管理器窗口中应该会出现一个简单的预置 CUDA 测试程序，包含一个源文件 kernel.cu，它包含一个简单的并行加法内核和测试代码。如果你想查看它是否可以编译和运行，请点击 IDE 顶部标记为本地 Windows 调试器的绿色右箭头。应该会弹出一个终端，显示内核的一些文本输出，然后立即关闭。

如果你从 Visual Studio 运行基于 Windows Terminal 的应用程序后遇到问题，尝试在主函数的末尾添加 getchar();，这将使终端保持打开状态，直到你按下键。 (或者，你还可以在程序末尾使用调试器断点。)

现在，让我们添加我们刚刚编写的 CUDA-C 程序。在解决方案资源管理器窗口中，右键单击 kernel.cu，然后点击 kernel.cu 上的移除。现在，右键单击项目名称，然后选择添加，然后选择现有项。现在我们将能够选择一个现有文件，所以找到 matrix_ker.cu 的路径并将其添加到项目中。点击 IDE 顶部标记为本地 Windows 调试器的绿色箭头，程序应该会编译并运行，再次在 Windows Terminal 中。所以，就是这样——我们只需这些步骤就可以在 Visual Studio 中设置和编译完整的 CUDA 程序了。

现在我们来看如何调试我们的 CUDA 内核。让我们首先在我们的代码中添加一个断点，在内核 matrix_mult_ker 的入口点，我们设置 row 和 col 的值。我们可以通过点击窗口左侧行号左侧的灰色列来添加这个断点；对于每个添加的断点，那里应该会出现一个红色圆点。（你可以忽略 Visual Studio 编辑器可能在你的代码下放置的任何红色波浪线；这是由于 CUDA 不是 Visual Studio 的 原生 语言）：

我们现在可以开始调试。从顶部菜单中选择 Nsight 下拉菜单，然后选择开始 CUDA 调试。这里可能有两个选项，开始 CUDA 调试（下一代）和开始 CUDA 调试（旧版）。哪个都无所谓，但您可能根据您的 GPU 遇到 Next-Gen 的问题；在这种情况下，请选择 Legacy。

您的程序应该启动，并且调试器应该在我们刚刚设置的内核断点处停止。让我们按F10来跳过这一行，现在看看row变量是否被正确设置。让我们查看变量探索器中的局部变量窗口：

我们可以看到，我们目前位于网格中第一个块中的第一个线程，通过检查threadIdx和blockIdx的值；row被设置为0，这确实对应于正确的值。现在，让我们检查不同线程的row值。为此，我们必须在 IDE 中切换线程焦点；我们通过点击 Nsight 下拉菜单上面的 Nsight，然后选择窗口|CUDA 调试焦点...来完成此操作。应该出现一个新菜单，允许您选择新的线程和块。在菜单中将线程从 0, 0, 0 更改为 1, 0, 0，然后点击确定：

当您再次检查变量时，应该看到row变量为这个线程设置了正确的值：

简而言之，这就是在 Visual Studio 中使用 Nsight 进行调试的方法。我们现在有了如何在 Windows 中的 Nsight/Visual Studio 中调试 CUDA 程序的基础，我们可以像在常规 Windows 程序中一样使用所有常规约定进行调试（设置断点、启动调试器、继续/恢复、跳过、进入和退出）。主要区别在于您必须知道如何在不同 CUDA 线程和块之间切换以检查变量，否则基本上是相同的。

在 Linux 中使用 Eclipse 的 Nsight

现在，我们将看到如何在 Linux 中使用 Nsight。您可以通过选择它或使用nsight命令从命令行运行它来打开 Nsight。Nsight IDE 将打开。从 IDE 的顶部，点击文件，然后从下拉菜单中选择新建...，然后从那里选择新建 CUDA C/C++项目。将出现一个新窗口，从这里选择 CUDA 运行时项目。给项目起一个合适的名字，然后点击下一步。您将需要提供进一步的设置选项，但默认设置对我们目前的目的来说已经足够好。（请注意，在第三和第四屏幕中，这里将位于源文件夹和项目路径的位置。）您将到达一个最终屏幕，在那里您可以按完成来创建项目：

最后，您将到达一个项目视图，其中包含您的新项目和一些占位符代码打开；截至 CUDA 9.2，这将是倒数第二个内核示例。

我们现在可以导入我们的代码。你可以只是使用 Nsight 中的编辑器删除默认源文件中的所有代码，并将其剪切粘贴进来，或者你可以手动从项目的源目录中删除文件，手动将 matrix_ker.cu 文件复制到源目录中，然后通过选择它并按 F5 来在 Nsight 中刷新源目录视图。你现在可以通过按 Ctrl + B 构建项目，并通过按 F11 运行它。程序输出应该出现在 IDE 的控制台子窗口中，如下所示：

我们现在可以在我们的 CUDA 代码中设置一个断点；让我们将其设置在内核入口点，此时行值被设置。我们在 Eclipse 编辑器中将光标置于该行上，然后按 Ctrl + Shift + B 来设置它。

我们现在可以通过按 F11（或点击错误图标）开始调试。程序应该在 main 函数的非常开始处暂停，因此按 F8 来 恢复 到第一个断点。你应该在 IDE 中看到我们的 CUDA 内核的第一行被高亮显示，并有一个箭头指向它；让我们通过按 F6 跳过当前行，这将确保行值已被设置。

现在，我们可以轻松地在我们的 CUDA 网格中的不同线程和块之间切换，以检查它们当前持有的值，如下所示：从 IDE 的顶部，点击窗口下拉菜单，然后点击显示视图，然后选择 CUDA。应该会打开一个包含当前正在运行的内核的窗口，从这里你可以看到该内核正在运行的块列表。

点击第一个，从这里你将能够看到块内运行的各个单独的线程：

现在，我们可以通过单击变量选项卡来查看第一个块中第一个线程对应的变量——在这里，行应该是 0，正如我们所预期的：

现在，我们可以通过再次转到 CUDA 选项卡，选择适当的线程，并切换回来来检查不同线程的值。让我们保持在同一个块中，但这次选择线程 (1, 0, 0)，再次检查行值：

我们看到行的值现在是 1，正如我们所预期的。

现在我们已经了解了如何在 Linux 中的 Nsight/Eclipse 上调试 CUDA 程序的基础，我们可以使用所有常规约定，就像在任何一个其他 IDE 中调试常规 Linux 程序一样（设置断点、启动调试器、继续/恢复、跳过、进入和退出）。主要区别在于，我们必须知道如何在不同 CUDA 线程和块之间切换以检查变量，否则基本上是相同的。

使用 Nsight 理解 CUDA 中的 warp lockstep 属性

我们现在将使用 Nsight 逐步执行一些代码，以帮助我们更好地理解一些 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存储库中的文件提供。）

如果我们从命令行编译并运行这个程序，我们可能会天真地期望看到来自偶数和奇数线程的字符串交错序列；或者它们可能会随机交错——因为所有线程都是并发运行的，并且大约在同一时间分支，这听起来是有道理的。

然而，每次我们运行这个程序时，我们总是得到这个输出：

所有对应于偶数线程的字符串首先打印出来，而所有奇数字符串随后打印。也许 Nsight 调试器可以提供一些线索；让我们像上一节那样将这个小程序导入到 Nsight 项目中，在内核中的第一个if语句处设置一个断点。然后我们将执行step over，这样调试器就会在第一个printf语句处停止。由于 Nsight 中的默认线程是（0,0,0），这应该满足了第一个if语句，所以它会卡在那里，直到调试器继续。

让我们切换到一个奇数线程，比如（1,0,0），看看它在我们的程序中的位置：

非常奇怪！线程（1,0,0）在执行过程中也处于与线程（0,0,0）相同的位置。实际上，如果我们检查这里所有的其他奇数线程，它们都会卡在相同的位置——在一个所有奇数线程都应该直接跳过的printf语句上。

发生了什么？这被称为** warp 步调一致属性。在 CUDA 架构中，warp** 是一个包含 32 个 "lanes" 的单元，我们的 GPU 在其中执行内核和网格，每个 lane 将执行一个线程。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 可执行程序进行分析：

我们看到这与我们最初在 第一章 中分析 Mandelbrot 算法的 Python cProfiler 模块的输出非常相似，为什么进行 GPU 编程？——但现在，这仅告诉我们所有已执行的 CUDA 操作。因此，当我们特别想要在 GPU 上进行优化时，我们可以使用这个工具，而不是关心主机上执行的任何 Python 或其他命令。 (如果我们添加命令行选项 --print-gpu-trace，我们还可以进一步分析每个单独的 CUDA 内核操作及其块和网格大小启动参数。)

让我们再来看一个技巧，帮助我们可视化程序中所有操作的执行时间；我们将使用nvprof来导出一个文件，然后可以被 NVIDIA Visual Profiler 读取，并以图形方式展示给我们。让我们用一个来自上一章的例子来演示，multi-kernel_streams.py（这个文件在仓库的5目录下有提供）。让我们回顾一下，这是我们介绍 CUDA 流概念的入门示例之一，CUDA 流允许我们并发执行和组织多个 GPU 操作。让我们使用以下命令将输出保存到以.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. 在 CUDA-C 中，cudaSetDevice命令的目的是什么？

7. 为什么在 CUDA-C 中，每次内核启动或内存复制后我们都必须使用cudaDeviceSynchronize？

第七章：使用 Scikit-CUDA 与 CUDA 库

在本章中，我们将游览三个旨在简化数值和科学计算的标准化 CUDA 库。我们将首先查看的是 cuBLAS，这是 NVIDIA 对 CUDA 的 基本线性代数子程序（BLAS）规范的实现。（cuBLAS 是 NVIDIA 对各种基于 CPU 的优化 BLAS 实现的回应，例如免费/开源的 OpenBLAS 或英特尔专有的数学内核库。）我们将查看的下一个库是 cuFFT，它可以在 GPU 上执行几乎所有的 快速傅里叶变换（FFT）变体。我们将特别探讨如何使用 cuFFT 进行图像处理中的滤波。然后我们将查看 cuSolver，它可以执行比 cuBLAS 中提供的更复杂的线性代数操作，例如 奇异值分解（SVD）或 Cholesky 分解。

到目前为止，我们主要处理的是一个充当我们进入 CUDA 的门户的单个 Python 模块——PyCUDA。虽然 PyCUDA 是一个非常强大且多功能的 Python 库，但其主要目的是提供进入程序、编译和启动 CUDA 内核的门户，而不是提供对 CUDA 库的接口。为此，幸运的是，有一个免费的 Python 模块可用，它提供了一个用户友好的包装器接口来访问这些库。这被称为 Scikit-CUDA。

即使你不需要了解 PyCUDA 或甚至理解 GPU 编程就能欣赏 Scikit-CUDA，但它与 PyCUDA 的兼容性非常方便；例如，Scikit-CUDA 可以轻松地与 PyCUDA 的 gpuarray 类一起操作，这允许你轻松地在我们的 CUDA 内核例程和 Scikit-CUDA 之间传递数据。此外，大多数例程也将与 PyCUDA 的流类一起工作，这将使我们能够正确地同步我们的自定义 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 用户。）或者，您可以直接从 PyPI 仓库使用pip install scikit-cuda来安装 Scikit-CUDA。

使用 cuBLAS 进行基本线性代数

我们将从这个章节开始，学习如何使用 Scikit-CUDA 的 cuBLAS 封装器。让我们花一点时间来讨论 BLAS。BLAS（基本线性代数子例程）是一个关于基本线性代数库的规范，它最初在 20 世纪 70 年代被标准化。BLAS 函数被分为几个类别，这些类别被称为级别。

Level 1 BLAS 函数包括仅对向量进行的操作——向量加法和缩放（也称为ax+y操作，或 AXPY），点积和范数。Level 2 BLAS 函数包括通用矩阵-向量操作（GEMV），例如向量的矩阵乘法，而 level 3 BLAS 函数包括“通用矩阵-矩阵”操作（GEMM），例如矩阵-矩阵乘法。最初，这些库在 20 世纪 70 年代完全用 FORTRAN 编写，因此您应该考虑到在用法和命名中可能存在一些看似过时的遗留问题，这些可能对今天的初学者来说显得繁琐。

cuBLAS 是 NVIDIA 自行实现的 BLAS 规范，当然是为了充分利用 GPU 的并行性而进行了优化。Scikit-CUDA 为 cuBLAS 提供了包装器，这些包装器与 PyCUDA gpuarray 对象以及 PyCUDA 流兼容。这意味着我们可以通过 PyCUDA 将这些函数与我们的自定义 CUDA-C 内核耦合和接口，同时还可以通过多个流同步这些操作。

cuBLAS 的 Level-1 AXPY

让我们从 cuBLAS 的基本 level-1 ax + y（或 AXPY）操作开始。让我们暂停一下，回顾一下线性代数，并思考这意味着什么。在这里，a 被认为是标量；也就是说，一个实数，例如 -10、0、1.345 或 100。x 和 y 被认为是某个向量空间中的向量，![img/47a6873c-3e1b-4b3c-95e8-d1a3a4f796eb.png]。这意味着 x 和 y 是由实数组成的 n-元组，所以在 ![img/d0e81dc7-0fa8-4bce-a264-941fee2e3ad7.png] 的情况下，这些可以是 [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 开头。

其他 level-1 cuBLAS 函数

让我们看看几个其他的 level-1 函数。我们不会深入探讨它们的操作，但步骤与我们刚刚覆盖的类似：创建 cuBLAS 上下文，使用适当的数组指针调用函数（这些指针通过 PyCUDA 的gpuarray的gpudata参数访问），并相应地设置步长。另一件需要注意的事情是，如果一个函数的输出是一个单值而不是数组（例如，点积函数），该函数将直接将此值输出到主机，而不是在从 GPU 拉取的内存数组中。（我们在这里只涵盖单精度实数版本，但其他数据类型的相应版本可以通过将 S 替换为适当的字母来使用。）

我们可以在两个单精度浮点 gpuarray，v_gpu 和 w_gpu 之间执行点积。再次，1s 是为了确保我们在这种计算中使用步长-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 中的 Level-2 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 中的 Level-3 GEMM 用于测量 GPU 性能

接下来，我们将探讨如何使用 CuBLAS 执行一个通用矩阵-矩阵乘法（GEMM）。实际上，我们将尝试做一些比我们在 cuBLAS 中看到的最后几个例子更有实用性的东西——我们将使用这个作为我们 GPU 性能的指标，以确定它每秒可以执行的浮点运算次数（FLOPS），这将有两个不同的值：单精度和双精度。使用 GEMM 是评估计算硬件性能的标准技术，因为它比使用纯时钟速度（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 - 1个加法，这意味着在A的m行中总共是km + (k-1)m个操作。B有n列，因此计算AB将总共是kmn + (k-1)mn = 2kmn - mn个操作。现在，我们使用alpha来缩放AB，这将涉及m**n个操作，因为这是矩阵AB的大小；同样，通过beta缩放C是另一个m**n操作。最后，我们将这两个结果矩阵相加，这又是另一个m个操作。这意味着在给定的 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）：

通常，计算矩阵-向量运算的计算复杂度为 O(N²)，其中 N 是向量的长度。然而，由于 DFT 矩阵中的对称性，这可以通过使用 FFT 总是减少到 O(N log N)。让我们看看我们如何使用 FFT 与 CuBLAS，然后我们将转向一个更有趣的例子。

简单的 1D FFT

让我们先看看如何使用 cuBLAS 来计算一个简单的 1D FFT。首先，我们将简要讨论 Scikit-CUDA 中的 cuFFT 接口。

这里有两个子模块，我们可以通过它们访问 cuFFT 库，分别是 cufft 和 fft。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)

我们现在将设置一个 cuFFT 计划用于前向 FFT 变换。这是一个 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

我们还没有完成。我们的卷积核将比我们的输入图像小得多，因此我们必须调整我们两个二维数组（卷积核和图像）的大小，使它们相等，并在它们之间执行逐点乘法。我们不仅要确保它们相等，还需要确保我们在数组上执行零填充，并且适当地居中卷积核。零填充意味着我们在图像的两侧添加一个零缓冲区，以防止环绕错误。如果我们使用 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 技巧来并排查看“之前”和“之后”的图像：

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 维的 dataset。我们首先将创建两个向量，它们在前面的权重很大，其他地方为 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]) ]

我们将随后添加 9,000 个额外的向量：其中 6,000 个与最初的两个向量相同，只是增加了一点点随机白噪声，剩下的 3,000 个则完全是随机白噪声：

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 库的包装器开始，这里我们需要记住许多细节，例如何时使用列主存储，或者一个输入数组是否会被就地覆盖。然后我们探讨了如何使用 Scikit-CUDA 的 cuFFT 执行一维和二维 FFT，以及如何创建一个简单的卷积滤波器。接着我们展示了如何将此应用于图像的简单高斯模糊效果。最后，我们探讨了如何使用 cuSolver 在 GPU 上执行奇异值分解（SVD），这通常是一个非常计算密集的操作，但它在 GPU 上并行化得相当好。我们本章的结尾是探讨如何使用 SVD 进行基本的 PCA。

问题

1. 假设你得到一份工作，将一些旧的遗留 FORTRAN BLAS 代码翻译成 CUDA。你打开一个文件，看到一个名为 SBLAH 的函数，另一个名为 ZBLEH。你能不查阅它们来告诉这两个函数使用的数据类型吗？

2. 你能否修改 cuBLAS 级别-2 GEMV 示例，使其通过直接将矩阵A复制到 GPU 上工作，而不需要在主机上转置以设置列向？

3. 使用 cuBLAS 32 位实数点积(cublasSdot)，通过一个行矩阵和一个步长为 1 的向量实现矩阵-向量乘法。

4. 使用cublasSdot实现矩阵-矩阵乘法。

5. 你能实现一种方法来精确测量性能测量示例中的 GEMM 操作吗？

6. 在 1D 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）。这将使我们能够以更自然的方式操作 CUDA C 数组，这种方式更接近 PyCUDA 的gpuarray和 STL 的 vector 容器。这将使我们免于不断使用指针，如mallocs和frees，这些在之前的 CUDA C 中困扰了我们。

本章我们将探讨以下主题：

• 理解种子在生成伪随机数列表中的作用

• 在 CUDA 内核中使用 cuRAND 设备函数生成随机数

• 理解蒙特卡洛积分的概念

• 在 Python 中使用基于字典的字符串格式化进行元编程

• 使用 CUDA Math API 设备函数库

• 理解什么是函子

• 在纯 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 中为 curand_uniform）是一个在给定范围内所有值等可能输出的函数：也就是说，对于 0 到 1 的均匀分布，有 10% 的概率值将落在 0 和 0.1 之间，或者落在 0.9 到 1 之间，或者落在彼此之间相隔 0.1 的任意两点之间。正态分布（在 cuRAND 中为 curand_normal）的值以特定的均值为中心，将根据分布的标准差分布，形成著名的钟形曲线。（在 cuRAND 中，curand_normal 的默认均值为 0，标准差为 1，因此对于其他值需要手动进行平移和缩放。）cuRAND 还支持另一个著名的分布，即泊松分布（curand_poisson），用于对随机事件随时间发生的情况进行建模。

在下一节中，我们将主要探讨如何在均匀分布的上下文中使用 cuRAND，因为它们适用于蒙特卡洛积分。对学习如何使用 cuRAND 的更多功能感兴趣的读者应查阅 NVIDIA 的官方文档。

使用蒙特卡洛估计 π

首先，我们将应用我们对 cuRAND 的新知识来估计著名的数学常数 π，也就是 Pi，它当然是一个无限不循环的无理数 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。然而，如果我们选择一个非常大的随机点值，请注意，我们将得到以下近似值：

这正是我们将如何估计 π！在我们能够得到 Pi 的合理估计之前，我们必须进行非常高的迭代次数，但请注意，这非常易于并行化：我们可以在不同的线程中检查“命中”，将总迭代次数分配给不同的线程。最终，我们只需将所有线程中的总“命中”数相加，以得到我们的估计。

我们现在可以开始编写一个程序来制作我们的蒙特卡洛估计。让我们首先导入我们将需要的常规 Python 模块，对于 PyCUDA 程序有一个来自 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。这修改了代码的编译方式，以便我们的代码可以正确地与 C++ 代码链接，这是 cuRAND 库所要求的。然后我们开始编写我们的内核并包含适当的头文件：

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;

让我们将全局线程 ID 存储在一个名为 tid 的整数中：

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 运行这个程序并检查它（此程序也作为monte_carlo_pi.py文件包含在此书存储库的Chapter08中）：

CUDA Math API

现在，我们将看看CUDA Math API。这是一个由与标准 C math.h库中类似的设备函数组成的库，可以从内核中的单个线程中调用。这里的一个区别是单精度和双精度浮点运算被重载，所以如果我们使用sin(x)，其中x是一个浮点数，那么正弦函数将输出一个 32 位浮点数，而如果x是一个 64 位双精度数，那么sin的输出也将是一个 64 位值（通常，这是 32 位函数的正确名称，但它在末尾有一个f，例如sinf）。还有额外的内建函数。内建函数不如原始函数精确，但速度更快，是集成在 NVIDIA CUDA 硬件中的；通常，它们的名称与原始函数相似，但前面有两个下划线——因此，内建、32 位正弦函数是__sinf。

定积分的简要回顾

现在，我们将使用 Python 中的面向对象编程来设置一个类，我们可以使用它来评估函数的定积分，使用蒙特卡洛方法。让我们停下来谈谈我们的意思：假设我们有一个数学函数（就像你在微积分课程中可能看到的那种），我们称之为f(x)。当我们在这两点a和b之间的笛卡尔平面上绘制这个函数时，它可能看起来像这样：

现在，让我们来回顾一下定积分的确切含义——让我们将此图中第一个灰色区域表示为 I，第二个灰色区域表示为 II，第三个灰色区域表示为 III。注意，这里的第二个灰色区域位于零以下。f 在这里的定积分，从 a 到 b，将是 I - II + III 的值，我们将这个数学上表示为 。一般来说，从 a 到 b 的定积分就是所有由 f 函数和 x 轴以及 y > 0 在 a 和 b 之间的“正”面积的总和，减去所有由 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 和 double 之间切换。让我们将字符串中的所有 float 引用更改为 %(precision)s——code_string="%(precision)s x, y; %(precision)s * z;"。现在我们可以设置一个合适的字典，将 %(presision)s 与 double 交换，即 code_dict = {'precision' : 'double'}，然后使用 code_double = code_string % code_dict 获取新的 double 字符串。让我们看一下：

现在，让我们暂时思考一下我们希望我们的新蒙特卡洛积分器如何工作。我们还将让它接受一个字符串，该字符串是一个使用 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;
}

现在，让我们思考一下这将如何工作——我们声明一个 32 或 64 位的浮点值y，调用math_function，然后返回y。math_function只有在它是某种作用于输入参数x并设置某些值到y的代码时才有意义，例如y = sin(x)。让我们记住这一点并继续。

现在，我们将开始编写我们的蒙特卡洛积分内核。让我们记住，我们必须使用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_lo 到 t_hi 的子积分值：

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 值。让我们还要确保在用户输入无效数据类型或 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)

我们现在准备尝试一下。让我们只设置一个具有默认值的类——这将把 y = sin(x) 从 0 到 π 进行积分。如果你记得微积分，sin(x) 的反导数是 -cos(x)，因此我们可以这样评估定积分：

因此，我们应该得到一个接近 2 的数值。让我们看看我们得到的是什么：

编写一些测试用例

现在，我们将最终看到如何通过 math_function 参数使用 CUDA Math API 为我们的类编写一些测试用例。如果你有任何 C/C++ 标准数学库的经验，这些将相当直接。同样，这些函数是重载的，这样我们就不必在单精度和双精度之间切换时更改任何名称。

我们已经看到了一个例子，即 y = sin(x)。让我们尝试一个稍微更有雄心的例子：

我们将从这个函数的 a=11.733 到 b=18.472 进行积分，然后检查我们的蒙特卡洛积分器的输出与另一个来源已知此积分值的比较。在这里，Mathematica 指出，这个定积分的值是 8.9999，因此我们将对此进行检查。

现在，让我们考虑如何表示这个函数：在这里，*log*指的是自然对数（也称为*ln*），在 Math API 中这仅仅是log(x)。我们已为平方设置了一个宏，因此我们可以将*sin²(x)*表示为_P2(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 中以更少的指针、mallocs 和 frees 依赖性进行编程。与 C++ 向量容器一样，Thrust 的向量容器自动处理元素的调整大小和连接，并且利用 C++ 析构函数的魔力，当 Thrust 向量对象超出作用域时，*freeing*也会自动处理。

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;

到目前为止，这可能看起来很 trivial。让我们在 GPU 上设置一个 Thrust 向量，然后从 v 中复制内容：

thrust::device_vector<int> v_gpu = v;

是的，就是这样——只有一行，我们就完成了。现在主机上的 v 的所有内容都将复制到设备上的 v_gpu！（如果这没有让你感到惊讶，请再次查看第六章，调试和性能分析您的 CUDA 代码，并思考在没有 Thrust 的情况下这需要多少行代码。）

让我们尝试在我们的新 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。）

让我们只设置第一个向量v，使用push_back：

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 积分示例（位于 monte_carlo_integrator.py 文件中的 __main__ 函数），使用 CUDA 的 instrinsic 函数。与之前相比，精度如何？

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、Microsoft 的 CNTK、Facebook 的 Caffe2 和 PyTorch，但从头开始实现一个是非常有教育意义的，这将使我们更深入地了解 DNN 所需的底层技术。这里有大量的材料要介绍，所以我们在简要介绍一些基本概念后，将直接进入正题。

在本章中，我们将探讨以下内容：

• 理解人工神经元（AN）是什么

• 理解在一个深度神经网络（DNN）中可以组合多少个 AN

• 在 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） 的基础知识。在机器学习中，我们的目标是收集一组具有特定标签类别或特征的数据，并使用这些示例来训练我们的系统以预测未来数据的值。我们称根据先前训练数据预测未来数据类别或标签的程序或函数为 分类器。

有许多类型的分类器，但在这里我们将专注于神经网络。神经网络背后的想法是它们（据说）以类似于人类大脑的方式工作，即它们通过使用一组 人工神经元（AN） 的集合来学习和分类数据，所有这些神经元都连接在一起形成一个特定的结构。然而，让我们暂时退后一步，看看单个 AN 是什么。在数学上，这只是一个从线性空间 R^n 到 R 的 仿射 函数，如下所示：

我们可以看到，这可以被描述为一个常数权重向量 w 与输入向量 x 之间的点积，并在末尾添加一个额外的偏置常数 b。 (再次强调，这个函数的唯一 输入 是 x；其他值都是常数！)

现在，单个 人工神经网络（AN） 个体来说相当无用（而且愚蠢），因为它们的 智能 只有在与大量其他 AN 协作时才会出现。我们的第一步是将一组 m 个类似的 AN 逐个堆叠起来，以形成一个我们称之为 密集层（DL） 的结构。这是密集的，因为每个神经元将处理来自 x 的每个单个输入值 - 每个 AN 将从 R^n 中接收一个数组或向量值，并输出一个 R 中的单个值。由于有 m 个神经元，这意味着我们可以说它们的输出总体上是在 R^m 空间中。我们会注意到，如果我们把每层的权重堆叠起来，形成一个 m x n 的权重矩阵，我们就可以通过矩阵乘法加上适当的偏置来计算每个神经元的输出：

现在，假设我们想要构建一个能够对 k 个不同的类别进行分类的神经网络分类器；我们可以创建一个新的额外的密集层，它接收来自先前密集层的 m 个值，并输出 k 个值。假设我们为每一层都找到了适当的权重和偏置值（这当然不是一件简单的事情），并且我们还在每一层之后设置了适当的 激活函数（我们将在后面定义），这将作为我们 k 个不同类别之间的分类器，根据最终层的输出给出 x 落入每个相应类别的概率。当然，我们在这里已经走得有点远了，但简而言之，这就是神经网络的工作原理。

现在，看起来我们只需将密集层连接成长链，就可以实现分类。这就是所谓的深度神经网络（DNN）。当我们有一个不直接连接到输入或输出的层时，它被称为隐藏层。DNN 的优势是额外的层允许神经网络捕捉到浅层神经网络无法捕捉到的数据的抽象和细微之处。

实现人工神经元的密集层

现在，让我们实现神经网络最重要的构建块，即 密集层。让我们首先声明一个 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（或任何其他激活函数）在神经网络中的隐藏层之间用作使整个神经网络作为非线性函数的手段；否则，整个神经网络将构成一个平凡的（且计算效率低下的）矩阵运算。（虽然还有许多其他非线性激活函数可以在层之间使用，但 ReLU 被发现是训练中特别有效的函数。）Sigmoid 被用作神经网络中用于标记的最终层，即可能为给定输入分配多个标签的层，而不是将输入分配给单个类别。

在我们甚至开始定义这个 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 类中，以便于使用，这将使我们在开始将这些密集层连接成一个完整的神经网络时生活得更加轻松。我们将称之为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 中存储为每个层的权重和输入/输出分配的内存。我们还将存储 DenseLayer 和 SoftmaxLayer 对象在列表 network 中，以及关于 NN 中每个层的信息在 network_summary 中。注意我们还可以在这里设置一些训练参数，包括 delta、用于梯度下降的流数量（我们稍后将看到），以及训练的轮数。

我们还可以在开始时看到一个名为 layers 的其他输入。在这里，我们可以通过描述每个层来指示神经网络的结构，构造函数将通过遍历 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 列表变量追加适当的对象和信息，以及适当地分配 gpuarray 对象到 network_mem 列表：

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)

现在，我们将从神经网络、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）的形式实现我们神经网络的完整训练方法。让我们逐字思考这意味着什么。批量意味着这个训练算法将一次处理一组训练样本，而不是同时处理所有样本，而随机表示每个批量是随机选择的。梯度意味着我们将使用微积分中的梯度——在这里，是损失函数上每个权重和偏置的导数集合。最后，下降意味着我们正在尝试减少损失函数——我们通过迭代地对权重和偏置进行细微的减法更改来实现这一点。

从微积分中记住，一个点的梯度始终指向最大增加的方向，其相反方向是最大减少的方向。由于我们想要一个减少，所以我们减去梯度。

现在，我们将实现 BSGD 作为我们SequentialNetwork类中的bsgd方法。让我们逐一过一下bsgd的输入参数：

• training将是一个二维 NumPy 数组，包含训练样本

• labels将是每个训练样本对应的神经网络最终层的期望输出

• 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

接下来，我们将遍历我们神经网络中的每个密集层，计算整个权重和偏置集的梯度。我们将这些权重和偏置的导数存储在展平（一维）数组中，这些数组将对应于我们的 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！

数据的预处理和标准化

在我们继续进行训练和测试我们全新的神经网络之前，我们需要暂时退后一步，谈谈数据预处理和数据标准化。神经网络对数值误差非常敏感，尤其是当输入具有较大尺度差异时。这可以通过正确预处理我们的训练数据来缓解；这意味着对于输入样本中的每个点，我们将计算所有样本中每个点的平均值和方差，然后在输入神经网络进行训练或推理（预测）之前，对每个样本中的每个点减去平均值并除以标准差。这种方法被称为标准化。让我们编写一个小的 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）文本文件，每行包含四个不同的数值（花瓣测量值），后面跟着花卉类型（在这里，有三个类别—Irissetosa，Irisversicolor和Irisvirginica）。现在，我们将设计一个小型 DNN 来根据这个集合对鸢尾花类型进行分类。

在我们继续之前，请下载爱丽丝数据集并将其放入您的当前工作目录。这个数据集可以从加州大学欧文分校机器学习存储库获得，网址如下：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 )

现在，让我们创建我们的神经网络。这将由四个密集层（两个隐藏层）和一个 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 分钟。

本章的代码存储在 deep_neural_network.py 文件中，位于本书 GitHub 仓库的相应目录下。

摘要

在本章中，我们首先给出了人工神经网络的定义，并展示了如何将单个人工神经网络（AN）组合成密集层，这些密集层最终组合成一个完整的深度神经网络。然后，我们在 CUDA-C 中实现了一个密集层，并创建了一个相应的 Python 包装类。我们还包含了在密集层的输出上添加 ReLU 和 sigmoid 层的功能。我们看到了使用 softmax 层的定义和动机，该层用于分类问题，并在 CUDA-C 和 Python 中实现了它。最后，我们实现了一个 Python 类，以便我们可以从之前的类构建一个顺序前馈 DNN；我们实现了一个交叉熵损失函数，然后在我们的 DNN 实现中使用这个损失函数来训练权重和偏差。最后，我们使用我们的实现在一个真实数据集上构建、训练和测试了一个 DNN。

由于我们能够编写自己的基于 GPU 的深度神经网络（DNN），我们现在对自己的 CUDA 编程能力非常有信心！接下来，我们将进入下一章的非常高级的内容，我们将探讨如何编写自己的接口来调用编译好的 CUDA 代码，以及一些关于 NVIDIA GPU 的非常技术性的细节。

问题

1. 假设你构建了一个 DNN，在训练后，它只产生了垃圾输出。经过检查，你发现所有的权重和偏置要么是巨大的数字，要么是 NaN。可能的问题是什么？

2. 列出一个小training_rate值可能存在的问题。

3. 列出一个大training_rate值可能存在的问题。

4. 假设我们想要训练一个 DNN，该 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 驱动程序 API 来结束这一章。然后我们可以使用驱动程序 API 中的适当函数来加载我们的 PTX 文件并启动内核。

本章的学习成果如下：

• 使用 Ctypes 模块启动编译后的（主机端）代码

• 使用 Ctypes 在 Python 中启动内核的主机端 CUDA C 包装器

• 如何将 CUDA C 模块编译成 PTX 文件

• 如何将 PTX 模块加载到 PyCUDA 中以启动预编译的内核

• 如何编写自己的自定义 Python 接口到 CUDA 驱动程序 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 中的函数，这些函数具有适当的 C 数据类型。

总是要确保将输入适当地类型转换为 Python 中使用 Ctypes 调用的任何函数的适当 C 数据类型（在 Ctypes 中，这些数据类型前面带有 c_：c_float、c_double、c_char、c_int 等）。

再次回顾曼德布罗特集（再次）

让我们回顾一下我们在 第一章 “为什么进行 GPU 编程？”和 第三章 “使用 PyCUDA 入门”中看到的曼德布罗特集。首先，我们将编写一个完整的 CUDA 内核，该内核将根据一组特定的参数计算曼德布罗特集，以及一个适当的主机端包装函数，我们可以在以后通过 Ctypes 接口。我们首先将这些函数写入一个单独的 CUDA-C .cu 源文件，然后使用 NVCC 编译器将此编译成 DLL 或 .so 二进制文件。最后，我们将编写一些 Python 代码，以便我们可以运行我们的二进制代码并显示曼德布罗特集。

现在，我们将应用我们对 Ctypes 的了解，从 Python 中启动预编译的 CUDA 内核，而不需要 PyCUDA 的任何帮助。这要求我们编写一个主机端 内核启动器 包装函数，我们可以直接调用它，该函数本身已被编译成动态库二进制文件，其中包含必要的 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_graph中。我们将使用max_iters指定每个点检查发散的迭代次数，通过提供其平方值upper_bound_squared来指定最大上限。（我们将在下一部分讨论使用平方的原因。）

我们将在一个一维网格/块结构上启动这个内核，每个线程对应于曼德布罗特集图形图像中的单个点。然后我们可以确定对应点的实/虚晶格值，如下所示：

    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坐标。由于晶格是一系列从小到大排序的实数值，我们将需要反转它们的顺序以获得适当的虚数。此外，请注意，我们将在这里使用普通的浮点数，而不是某种结构或对象来表示复数值。由于每个复数都有实部和虚部，我们在这里将需要使用两个浮点数来存储对应于这个线程晶格点的复数（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;
}

然而，我们还没有完全完成。我们需要为 Linux 编写一个仅包含extern "C"的主机端包装函数，而在 Windows 的情况下，则需要extern "C" __declspec(dllexport)。（与编译后的 CUDA 内核相比，如果我们想在 Windows 中从 Ctypes 访问主机端函数，这个额外的关键字是必要的。）我们放入这个函数的参数将直接对应于进入内核的参数，除了这些将存储在主机上：

extern "C" __declspec(dllexport) void launch_mandelbrot(float * lattice,  float * mandelbrot_graph, int max_iters, float upper_bound, int lattice_size)
{

现在，我们首先要做的是使用cudaMalloc为存储晶格和输出在 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。我们通过使用适当的 Ctypes 数据类型（c_float、c_int）以及 Ctypes 的 POINTER 类，通过 argtypes 参数来设置这些：

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)

现在，我们已经准备好计算曼德布罗特图了。请注意，我们可以通过使用它们的 ctypes.data_as 方法以及相应的类型将 NumPy 数组传递给 C。完成此操作后，我们可以返回输出；即，以二维 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()

我们现在尝试运行这个程序。你应该得到一个输出，其外观与第一章“为什么进行 GPU 编程？”和第三章“使用 PyCUDA 入门”中的曼德布罗特图完全一样：

这个 Python 示例的代码也作为 GitHub 仓库中的文件 mandelbrot_ctypes.py 提供。

编译和启动纯 PTX 代码

我们刚刚看到了如何从 Ctypes 调用一个纯 C 函数。在某种程度上，这可能看起来有点不太优雅，因为我们的二进制文件必须包含主机代码以及编译后的 GPU 代码，这可能会显得有些繁琐。我们能否只使用纯编译后的 GPU 代码，然后适当地将其加载到 GPU 上，而不必每次都编写 C 包装器？幸运的是，我们可以。

NVCC 编译器将 CUDA-C 编译成 PTX（并行线程执行），这是一种与 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 驱动 API 编写包装器

我们现在将探讨如何使用 Ctypes 为一些预包装的二进制 CUDA 库函数编写自己的包装器。特别是，我们将编写 CUDA 驱动 API 的包装器，这将使我们能够执行所有必要的操作，以实现基本的 GPU 使用——包括 GPU 初始化、内存分配/传输/释放、内核启动以及上下文创建/同步/销毁。这是一项非常强大的知识；它将使我们能够使用 GPU，而无需通过 PyCUDA，也无需编写任何繁琐的主机端 C 函数包装器。

我们现在将编写一个小的模块，它将作为CUDA 驱动 API的包装库。让我们先谈谈这意味着什么。驱动 API 与CUDA 运行时 API略有不同，后者是我们在这篇 CUDA-C 文本中一直在使用的，后者稍微技术性更强。驱动 API 旨在与常规的 C/C++编译器一起使用，而不是与 NVCC 一起使用，有一些不同的约定，例如使用cuLaunchKernel函数来启动内核，而不是使用<<< gridsize, blocksize >>>括号表示法。这将使我们能够直接访问从 PTX 文件启动内核所需的函数，使用 Ctypes。

让我们从将所有 Ctypes 导入模块命名空间开始编写这个模块，然后导入 sys 模块。我们将通过检查系统的 OS（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 上运行的进程类似。由于这是通过运行时 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++中，您始终可以使用void *类型（在 Ctypes 中为c_void_p）来指向任何任意的数据或变量——甚至定义可能不可用的结构和对象。

下一个函数是cuModuleLoad，它将为我们加载一个 PTX 模块文件。第一个参数是一个通过引用的 CUmodule（再次，我们在这里可以使用c_void_p），第二个参数是文件名，它将是一个典型的以 null 结尾的 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 驱动 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文件中。我们现在已经完成了我们的驱动 API 包装模块！接下来，我们将看看如何仅使用我们的模块和我们的 Mandelbrot PTX 来加载一个 PTX 模块并启动一个内核。

此示例也作为本书 GitHub 仓库中的cuda_driver.py文件提供。

使用 CUDA 驱动 API

现在，我们将翻译我们的 Mandelbrot 生成程序，以便我们可以使用我们的包装库。让我们从适当的导入语句开始；注意我们如何将所有的包装器加载到当前命名空间中：

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 驱动 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：这是 C 编程中引用操作符（&）的 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 中它是如何工作的。

我们将输入参数以 void * 值的数组形式表示在 kernelParams 中，这些值本身是指向我们想要插入内核的输入的指针。在我们的 Mandelbrot 内核的情况下，它看起来像这样：

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()

现在尝试运行这个函数，以确保它产生与我们所写的其他 Mandelbrot 程序相同的输出。

恭喜你——你刚刚编写了一个直接访问低级 CUDA 驱动 API 的接口，并成功使用它启动了一个内核！

此程序也作为 mandelbrot_driver.py 文件包含在此书 GitHub 仓库的目录下。

概述

我们在本章开始时简要概述了 Python Ctypes 库，该库用于直接与编译的二进制代码接口，特别是用 C/C++ 编写的动态库。然后我们探讨了如何使用 CUDA-C 编写一个包装器来启动 CUDA 内核，然后使用它通过编写 Ctypes 函数的接口间接从 Python 启动 CUDA 内核。然后我们学习了如何将 CUDA 内核编译成 PTX 模块二进制文件，这可以被视为 DLL，但包含 CUDA 内核函数，并看到了如何使用 PyCUDA 加载 PTX 文件并启动预编译的内核。最后，我们编写了一组 CUDA 驱动 API 的 Ctypes 包装器，并看到了如何使用这些包装器执行基本的 GPU 操作，包括从 PTX 文件启动预编译的内核到 GPU 上。

现在我们将进入本书可能最具技术性的章节：第十一章，CUDA 中的性能优化。在这一章中，我们将了解一些关于 NVIDIA GPU 的技术细节，这将帮助我们提高应用程序的性能水平。

问题

1. 假设你使用nvcc编译一个包含主机和内核代码的单一.cu文件，将其编译成 EXE 文件，同时也编译成 PTX 文件。哪个文件将包含主机函数，哪个文件将包含 GPU 代码？

2. 为什么在使用 CUDA 驱动程序 API 时，我们必须销毁上下文？

3. 在本章的开头，当我们第一次看到如何使用 Ctypes 时，请注意，在调用printf之前，我们必须将浮点值 3.14 强制转换为 Ctypes 的c_double对象，这样它才能正常工作。然而，在本章中我们可以看到许多不需要强制转换为 Ctypes 就能正常工作的案例。你认为为什么printf在这里是个例外？

4. 假设你想向我们的 Python CUDA 驱动程序接口模块添加功能以支持 CUDA 流。你将如何用 Ctypes 表示单个流对象？

5. 为什么我们在mandelbrot.cu中的函数使用extern "C"？

6. 再看看mandelbrot_driver.py。为什么我们在 GPU 内存分配和主机/GPU 内存传输之后，以及仅在单个内核调用之后才使用cuCtxSynchronize函数？

第十一章：CUDA 中的性能优化

在本章的最后，我们将介绍一些相当高级的 CUDA 特性，我们可以使用它们进行底层性能优化。我们将从学习动态并行开始，这允许内核在 GPU 上启动和管理其他内核，并看看我们如何使用它来实现直接在 GPU 上的快速排序。我们将学习关于矢量化内存访问的内容，这可以在从 GPU 的全局内存读取时提高内存访问速度。然后我们将看看如何使用 CUDA 原子操作，这些是线程安全的函数，可以在没有线程同步或互斥锁的情况下操作共享数据。我们将学习关于 Warps 的内容，它们是由 32 个或更少的线程组成的根本块，线程可以直接读取或写入彼此的变量，然后简要地进入 PTX 汇编的世界。我们将通过在 CUDA-C 代码中直接编写一些基本的 PTX 汇编内联来实现这一点，而 CUDA-C 代码本身将内联在我们的 Python 代码中！最后，我们将把这些小的底层调整集中到一个最终的例子中，我们将应用它们来制作一个极快的求和内核，并将其与 PyCUDA 的求和进行比较。

本章的学习成果如下：

• CUDA 中的动态并行

• 使用动态并行在 GPU 上实现快速排序

• 使用矢量化类型来加速设备内存访问

• 使用线程安全的 CUDA 原子操作

• 基本 PTX 汇编

• 将所有这些概念应用于编写性能优化的求和内核

动态并行

首先，我们将看看动态并行，这是 CUDA 中的一个特性，允许内核在没有主机交互或输入的情况下启动和管理其他内核。这也使得许多通常在主机端可用的 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²)。快速排序通过在未排序数组中选择一个任意点称为pivot，然后将数组划分为一个左数组（包含所有小于枢轴的点），一个右数组（包含所有等于或大于枢轴的点），枢轴位于两个数组之间来执行。如果现在一个或两个数组长度大于 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);

}'''

注意我们如何必须使用解引用运算符*来设置向量化变量，以及我们如何必须通过引用跳转到下一个地址（&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 为每个线程的 add_out 加 1，这将给出线程的总数：

 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 上的线程，这限制了 shuffling 操作只能用于大小为 32 或更小的线程组。另一个陷阱是，我们只能使用 32 位或更小的数据类型。这意味着我们无法在 Warp 中 shuffle 64 位 long long 整数或 double 浮点值。

只有 32 位（或更小）的数据类型可以与 CUDA Warp shuffling 一起使用！这意味着虽然我们可以使用整数、浮点数和字符，但不能使用双精度或长长整数！

在我们继续任何编码之前，简要回顾一下 CUDA Warps。（在继续之前，你可能希望回顾第六章中名为The warp lockstep property的部分，即调试和性能分析您的 CUDA 代码。）CUDA 中的Warp是 CUDA 中最小的执行单元，由 32 个或更少的线程组成，在恰好 32 个 GPU 核心上运行。正如 Grid 由块组成一样，块同样由一个或多个 Warp 组成，具体取决于块使用的线程数——如果一个块由 32 个线程组成，那么它将使用一个 Warp，如果它使用 96 个线程，它将包含三个 Warp。即使 Warp 的大小小于 32，它也被视为一个完整的 Warp：这意味着只有一个线程的块将使用 32 个核心。这也意味着一个包含 33 个线程的块将包含两个 Warp 和 31 个核心。

为了记住我们在第六章中看到的内容，即调试和性能分析您的 CUDA 代码，Warp 有一个被称为锁步属性的特性。这意味着 Warp 中的每个线程将迭代每条指令，与 Warp 中的每个其他线程完美并行。也就是说，单个 Warp 中的每个线程将同时执行相同的精确指令，忽略任何不适用于特定线程的指令——这就是为什么在单个 Warp 中尽可能避免线程之间的任何分歧。NVIDIA 称这种执行模型为单指令多线程，或SIMT。到现在为止，你应该明白为什么我们一直在文本中始终如一地使用 32 线程的块！

在我们继续之前，我们需要学习一个新术语——Warp 中的一个lane是 warp 中特定线程的唯一标识符，其值将在 0 到 31 之间。有时，这也被称为Lane ID。

让我们从简单的例子开始：我们将使用__shfl_xor命令在所有偶数和奇数编号的 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 的值作为输入。它对当前线程的二进制通道 ID 与 1 进行 XOR 操作，这将要么是其左邻居（如果线程的通道的最不重要位是二进制的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 混洗的例子——我们将实现一个操作，用于在 warp 中的所有线程上对单个局部变量求和。让我们回顾一下来自第四章，内核、线程、块和网格中的朴素并行求和算法，这个算法非常快，但它做出了一个朴素的假设，即我们拥有的处理器数量和我们拥有的数据块数量一样多——在我们生活中，这实际上是非常少见的，假设我们正在处理大小为 32 或更小的数组。我们将使用 __shfl_down 函数在一个 warp 中实现这一点。__shfl_down 函数接受第一个参数中的线程变量，通过第二个参数中指示的步数在线程之间移位一个变量，而第三个参数将指示 warp 的总大小。

现在我们立即实现这个操作。再次提醒，如果你不熟悉朴素并行求和或者不记得为什么它应该有效，请回顾第四章，内核、线程、块和网格。我们将使用 __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()) )

这将给出以下输出：

本节中的示例也作为 shfl_sum.py 和 shfl_xor.py 文件包含在本书的 GitHub 仓库中的 Chapter11 目录下。

内联 PTX 汇编

现在，我们将探讨编写 PTX（并行线程执行）汇编语言，这是一种在所有 Nvidia GPU 上工作的伪汇编语言，它反过来由即时（JIT）编译器编译成特定 GPU 的实际机器代码。虽然这显然不是日常使用，但如果需要，它将让我们在比 C 语言更低级别上工作。一个特定的用例是，你可以轻松地反汇编 CUDA 二进制文件（主机端可执行文件/库或 CUDA .cubin 二进制文件）并检查其 PTX 代码，如果否则没有源代码。这可以通过 Windows 和 Linux 中的cuobjdump.exe -ptx cuda_binary命令来完成。

如前所述，我们将在 CUDA-C 中仅涵盖 PTX（并行线程执行）的一些基本用法，CUDA-C 具有特定的语法和用法，这与在 GCC 中使用内联主机端汇编语言类似。让我们开始编写我们的代码——我们将进行导入并开始编写我们的 GPU 代码：

from __future__ import division
import numpy as np
from pycuda.compiler import SourceModule
import pycuda.autoinit
from pycuda import gpuarray

PtxCode='''

我们将通过将代码写入单独的设备函数来在这里进行几个迷你实验。让我们从一个简单的函数开始，该函数将输入变量设置为 0。（在 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使用与字符串这里的0th参数相对应的寄存器，我们用逗号将其与mov的下一个输入分开，逗号是常数0。然后我们像在 C 语言中一样以分号结束汇编语句行，并用引号关闭这段汇编代码。然后我们必须使用冒号（而不是逗号！）来指示我们想要在代码中使用哪些变量。"=r"意味着两件事：等号=将指示nvcc该寄存器将被写入作为输出，而r表示这应该被处理为 32 位整数数据类型。然后我们将要由汇编器处理的变量放在括号中，然后像任何 C 函数一样关闭asm。

所有这些都是为了设置单个变量的值为 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 来表示我们只从它读取。此外，注意我们如何使用冒号来区分 nvcc 中的 write 寄存器和 only read 寄存器。

在我们继续之前，让我们看看另一个简单的例子，即类似于 C 中的 ++ 操作符的函数，它将整数增加 1。

__device__ void plusplus(int &x)
{
 asm("add.s32 %0, %0, 1;" : "+r"(x));
}

首先，请注意我们使用“0th”参数作为输出和第一个输入。接下来，请注意我们使用 +r 而不是 =r——+ 告诉 nvcc，在这个指令中这个寄存器将被读取和写入。

现在，我们不会比这更复杂，因为即使在汇编语言中编写一个简单的 if 语句也是相当复杂的。然而，让我们看看一些更有用的例子，这些例子在使用 CUDA Warps 时会很有帮助。让我们从一个小的函数开始，它将给我们当前线程的通道 ID；这特别有用，实际上比使用 CUDA-C 做这件事要简单得多，因为通道 ID 实际上存储在一个称为 %laneid 的特殊寄存器中，我们无法在纯 C 中访问它。（注意我们在代码中使用两个 % 符号，这将指示 nvcc 直接在汇编代码中使用 % 来引用 %laneid，而不是将其解释为 asm 命令的参数。）

__device__ int laneid()
{
 int id; 
 asm("mov.u32 %0, %%laneid; " : "=r"(id)); 
 return id;
}

现在，让我们再写两个有用的函数来处理 CUDA Warps。记住，你只能通过 shuffle 命令传递一个 32 位变量。这意味着为了在 warp 中传递一个 64 位变量，我们必须将其分成两个 32 位变量，分别将这两个变量 shuffle 到另一个线程，然后将这些 32 位值重新组合成原始的 64 位变量。我们可以使用 mov.b64 命令来处理将 64 位双精度浮点数拆分成两个 32 位整数的情况——注意我们必须使用 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 自身的 sum 函数在 gpuarray 对象上的性能。

让我们开始导入所有必要的库，然后从与上一节类似的 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，它应该初始化为 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 变量的 double2 中求和到一个新的双精度变量 sum_val 中，这将跟踪这个线程的所有求和。我们将创建两个 32 位整数 s1 和 s2，我们将使用它们来拆分这个值并与 warp 混洗共享，然后创建一个 temp 变量来存储我们从这个 warp 中的其他线程接收到的重构值：

 double sum_val = vals.x + vals.y;

 double temp;

 int s1, s2;

现在，让我们再次使用 Naive Parallel sum 在 warp 中进行求和，这相当于在 warp 中对 32 位整数进行求和，但我们将在每个迭代中使用我们的 split64 和 combine64 PTX 函数在 sum_val 和 temp 上：

 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 线程使用线程安全的 atomicAdd 将它们的结束值添加到 out 中：

 if (id == 0)
     atomicAdd(out, sum_val);

}'''

我们现在将使用 timeit 操作编写测试代码，以测量我们的内核和 PyCUDA 的 sum 函数在 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 的编译而计时：

所以，虽然求和通常相当无聊，但我们却可以因为巧妙地使用硬件技巧，能显著加快这种平淡无奇且微不足道的算法而感到兴奋。

这个例子作为performance_sum_ker.py文件包含在这本书的 GitHub 仓库的Chapter11目录下。

摘要

我们从学习动态并行性开始这一章，这是一种允许我们从其他内核直接在 GPU 上启动和管理内核的范例。我们看到了如何直接在 GPU 上实现快速排序算法。然后我们学习了 CUDA 中的矢量化数据类型，并看到了如何使用这些数据类型来加速从全局设备内存中的内存读取。然后我们学习了 CUDA Warps，它们是 GPU 上 32 个线程或更少的单元，我们看到了单个 Warp 内的线程如何使用 Warp Shuffle 直接读取和写入彼此的寄存器。然后我们探讨了如何用 PTX 汇编编写一些基本操作，包括导入操作，如确定车道 ID 和将 64 位变量分割成两个 32 位变量。最后，我们通过编写一个新的用于双精度数组性能优化的求和内核来结束这一章，我们应用了本章学到的几乎所有技巧。我们看到这实际上比标准 PyCUDA 对长度为 500,000 的双精度数组求和要快。

我们已经完成了这本书的所有技术章节！你应该为自己感到骄傲，因为你现在肯定是一个熟练的 GPU 程序员，有很多技巧在你的袖子里。我们现在将开始最后一章，我们将简要地浏览一下你可以采取的不同路径来应用和扩展你的 GPU 编程知识。

问题

1. 在原子操作示例中，尝试在内核启动之前将网格大小从 1 更改为 2，同时保持总块大小为 100。如果这导致add_out（除了 200 以外的任何值）的输出不正确，那么为什么是错误的，考虑到atomicExch是线程安全的？

2. 在原子操作示例中，尝试移除__syncthreads，然后使用网格大小为 1 和块大小为 100 的原始参数运行内核。如果这导致add_out（除了 100 以外的任何值）的输出不正确，那么为什么是错误的，考虑到atomicExch是线程安全的？

3. 为什么我们不需要在大小为 32 或更小的块上使用__syncthreads来同步？

4. 我们看到sum_ker大约比 PyCUDA 对长度为 640,000（10000*2*32）的随机值数组求和快五倍。如果你尝试在这个数字的末尾添加一个零（即乘以 10），你会注意到性能下降到sum_ker只比 PyCUDA 的求和快 1.5 倍。如果你在这个数字的末尾再添加一个零，你会注意到sum_ker只比 PyCUDA 的求和快 75%。你认为这是为什么？我们如何改进sum_ker以使其在更大的数组上更快？

5. 哪个算法执行更多的加法操作（包括对 C +运算符和 atomicSum 的调用）：sum_ker还是 PyCUDA 的sum？

第十二章：从哪里开始

这本书就像一次大胆的山地徒步旅行一样，是一段旅程……但现在，我们终于到达了旅行的终点。我们现在站在入门级 GPU 编程的山顶上，我们自豪地回望我们的故乡——串行编程城，当我们想到我们旧的一维编程传统是多么的纯真时，我们会微笑。我们勇敢地克服了许多陷阱和危险才到达这个点，我们甚至可能犯了一些错误，比如在 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 因为过多的工作而饱和，而其他 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 等语言一起使用，允许你编程连接到同一网络的多台计算机。

关于使用 CUDA 与 MPI 的更多信息，请参阅此处：devblogs.nvidia.com/introduction-cuda-aware-mpi/。

OpenCL 和 PyOpenCL

CUDA 不是唯一可以用来编程 GPU 的语言。CUDA 最主要的竞争对手被称为开放计算语言，或 OpenCL。CUDA 是一个封闭且专有的系统，仅能在 NVIDIA 硬件上运行，而 OpenCL 是一个由非营利组织 Khronos Group 开发和支持的开放标准。OpenCL 可以用来编程不仅限于 NVIDIA GPU，还包括 AMD Radeon GPU 和英特尔 HD GPU——大多数主要科技公司都承诺在其产品中支持 OpenCL。此外，PyCUDA 的作者、UIUC 的安德烈亚斯·克洛克纳教授还编写了另一个优秀的（且免费）Python 库，名为 PyOpenCL，它提供了与 PyCUDA 几乎相同的语法和概念，为 OpenCL 提供了一个同样用户友好的接口。

关于 OpenCL 的信息由 NVIDIA 提供：developer.nvidia.com/opencl.

关于免费 PyOpenCL 库的信息可以从安德烈亚斯·克洛克纳的网站获取：

mathema.tician.de/software/pyopencl/.

图形

显然，GPU 中的 G 代表 图形，在这本书中我们并没有看到很多关于它的内容。尽管机器学习应用现在是 NVIDIA 的主要业务，但一切始于制作漂亮的图形。我们将在此提供一些资源，无论您是想开发视频游戏引擎、渲染 CGI 电影，还是开发 CAD 软件，都可以从这里开始。CUDA 实际上可以与图形应用结合使用，实际上它被用于 Adobe 的 Photoshop 和 After Effects 等专业软件，以及许多最近的游戏，如 Mafia 和 Just Cause 系列。我们将简要介绍一些您可能考虑从这里开始的主要 API。

OpenGL

开放图形语言，或 OpenGL，是一个自 90 年代初就存在的行业开放标准。虽然它在某些方面显示出其年代感，但它是一个稳定的 API，得到了广泛的支持，如果您编写了一个利用这个 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 可以被视为 Khronos Group 开发的 DirectX 12 的开放等效版本，它是 OpenGL 的下一代继任者。除了 Windows 外，Vulkan 还支持 macOS、Linux，以及索尼 PlayStation 4、任天堂 Switch 和 Xbox One 游戏机。Vulkan 具有与 DirectX 12 相似的功能，如准 GPGPU 编程。Vulkan 正在为 DirectX 12 提供一些严重的竞争，例如 2016 年的 DOOM 重制游戏。

《Vulkan 入门指南》可在 Khronos Group 这里获取：www.khronos.org/blog/beginners-guide-to-vulkan.

机器学习和计算机视觉

当然，本章中不容忽视的是机器学习和其孪生兄弟计算机视觉。不言而喻，机器学习（尤其是深度神经网络和卷积神经网络等子领域）是如今支撑着 NVIDIA 首席执行官黄仁勋的屋顶的东西。（好吧，我们承认这是本世纪最夸张的夸张……）如果你需要提醒为什么 GPU 在这个领域如此适用和有用，请再次查看第九章，深度神经网络的实现。大量的并行计算和数学运算，以及用户友好的数学库，使得 NVIDIA GPU 成为机器学习行业的硬件支柱。

基础知识

虽然你现在已经了解了底层 GPU 编程的许多复杂性，但你无法立即将这些知识应用到机器学习上。如果你在这个领域没有基本技能，比如如何对数据集进行基本的统计分析，你真的应该停下来熟悉它们。斯坦福大学教授、谷歌大脑的创始人 Andrew Ng 提供了许多免费材料，这些材料可以在网上和 YouTube 上找到。Ng 教授的工作通常被认为是机器学习教育材料的黄金标准。

Ng 教授在网上提供了一门免费的机器学习入门课程：www.ml-class.org.

cuDNN

NVIDIA 为深度神经网络原语提供了一个优化的 GPU 库，称为 cuDNN。这些原语包括前向传播、卷积、反向传播、激活函数（如 sigmoid、ReLU 和 tanh）和梯度下降等操作。cuDNN 是大多数主流深度神经网络框架（如 Tensorflow）用作 NVIDIA GPU 后端的东西。这是由 NVIDIA 免费提供的，但必须从 CUDA 工具包中单独下载。

关于 cuDNN 的更多信息在这里：developer.nvidia.com/cudnn。

Tensorflow 和 Keras

Tensorflow 当然是谷歌知名的神经网络框架。这是一个免费且开源的框架，可以用 Python 和 C++ 使用，自 2015 年以来对公众开放。

在这里可以找到 Tensorflow 的教程：www.tensorflow.org/tutorials/。

Keras 是一个更高级的库，它为 Tensorflow 提供了一个更用户友好的接口，最初由 Google Brain 的 Francois Chollet 编写。读者实际上可以考虑先从 Keras 开始，然后再转向 Tensorflow。

关于 Keras 的信息在这里：keras.io/。

Chainer

Chainer 是由日本东京大学的博士生 Seiya Tokui 开发的另一个神经网络 API。虽然它的主流程度不如 Tensorflow，但由于其惊人的速度和效率，它非常受尊重。此外，读者可能会对 Chainer 特别感兴趣，因为它是最初使用 PyCUDA 开发的。（后来改为使用 CuPy，这是 PyCUDA 的一个分支，旨在提供一个与 NumPy 更相似的接口。）

关于 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. 尝试找到至少一种可以用来编程本章未提及的 GPU 的编程语言或 API。

3. 查找谷歌的新 Tensor Processing Unit (TPU)芯片。这些芯片与 GPU 有何不同？

4. 你认为使用 Wi-Fi 还是有线以太网电缆将计算机连接成集群更好？

第十三章：评估

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

1. 前两个 for 循环遍历每个像素，其输出对彼此是恒定的；因此，我们可以在这两个 for 循环上并行化。第三个 for 循环计算特定像素的最终值，这是本质上是递归的。

2. Amdahl 定律没有考虑到在 GPU 和主机之间传输内存所需的时间。

3. 512 x 512 等于 262,144 像素。这意味着第一个 GPU 只能一次计算一半像素的输出，而第二个 GPU 可以一次计算所有像素；这意味着第二个 GPU 在这里将比第一个快大约两倍。第三个 GPU 有足够的内核一次计算所有像素，但正如我们在问题 1 中看到的，额外的内核在这里对我们没有帮助。因此，对于这个问题，第二个和第三个 GPU 的速度将一样快。

4. 关于 Amdahl 定律，将某段代码泛化为可并行化的代码有一个问题是，这假设如果处理器数量 N 非常大，这段代码的计算时间将接近 0。正如我们从最后一个问题中看到的，这并不成立。

5. 首先，一致地使用 time 可能会很麻烦，而且可能无法准确找到程序的性能瓶颈。其次，分析器可以从 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 与 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. cuCtxDestroy。

3. 带有任意输入参数的printf。（尝试查找printf原型。）

4. 使用 Ctypes 的c_void_p对象。

5. 这将允许我们从 Ctypes 链接到具有其原始名称的函数。

6. 设备内存分配和设备/主机之间的 memcopies 由 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. 以太网。