GPGPU Achitectures阅读笔记2

Chapter 2 Programming Model

本章的目标是提供足够的关于如何为非图形计算编程 GPU 的上下文，以便那些之前没有 GPU 经验的人可以遵循后面章节的讨论。 我们在这里关注基本材料，将更深入的报道留给其他参考文献（例如，[Kirk 和 Wen-Mei，2016]）。 存在许多可用于架构研究的 GPU 计算基准套件。 了解如何对 GPU 进行编程与对 GPU 计算感兴趣的计算机架构师相关，以便更好地了解硬件/软件接口，但如果您想探索对硬件/软件接口进行更改作为研究的一部分，它就变得必不可少。 在后一种情况下，现有的基准可能不存在，因此可能需要通过修改现有 GPU 计算应用程序的源代码来创建。 例如，探索在 GPU 上引入事务内存 ™ 的研究需要这样做，因为当前的 GPU 不支持 TM（见第 5.3 节）

现代 GPU 使用广泛的 SIMD 硬件来利用 GPU 应用程序中的数据级并行。 GPU 计算 API（例如 CUDA 和 OpenCL）不是直接向程序员公开这种 SIMD 硬件，而是具有类似 MIMD 的编程模型，允许程序员在 GPU 上启动大量标量线程。 这些标量线程中的每一个都可以遵循其独特的执行路径，并且可以访问任意内存位置。 在运行时，GPU 硬件在 SIMD 硬件上同步执行称为扭曲（或 AMD 术语中的波前）的标量线程组，以利用它们的规律性和空间局部性。 这种执行模型称为单指令多线程 (SIMT) [Lindholm 等，2008a，Nickolls 和 Reusch，1993]。

本章的其余部分在此讨论的基础上进行了扩展，并按如下方式组织。 在第 2.1 节中，我们探讨了最近 GPU 编程模型使用的概念执行模型，并简要总结了过去十年发布的典型 GPU 的执行模型。 在 2.2 节中，我们探讨了 GPU 计算应用程序的编译过程，并简要介绍了 GPU 指令集架构。

2.1 EXECUTION MODEL

GPU 计算应用程序开始在 CPU 上执行。 对于离散 GPU，应用程序的 CPU 部分通常会分配内存用于 GPU 上的计算，然后启动将输入数据传输到 GPU 内存，最后在 GPU 上启动计算内核。 对于集成 GPU，只需要最后一步。 计算内核由（通常）数千个线程组成。 每个线程执行相同的程序，但根据计算结果可能遵循不同的控制流通过该程序。 下面我们使用用 CUDA 编写的特定代码示例详细考虑此流程。 在下一节中，我们将查看程序集级别的执行模型。 我们的讨论不会停留在 GPU 编程模型的性能方面。 然而，Seo 等人的一项有趣观察。 [2011] 在 OpenCL（一种类似于 CUDA 的编程模型，可以编译为多种架构）的上下文中，针对一种架构（例如 GPU）精心优化的代码可能在另一种架构（例如 CPU）上表现不佳。

图 2.1 提供了 C 代码，用于 CPU 实现众所周知的操作单精度标量值 A 乘以向量值 X 加上向量值 Y，称为 SAXPY。 SAXPY 是著名的基本线性代数软件 (BLAS) 库 [Lawson et al., 1979] 的一部分，可用于实现更高级别的矩阵运算，例如高斯消元 [McCool等，2012]。 鉴于其简单性和实用性，在教授计算机体系结构时经常将其用作示例 [Hennessy and Patterson, 2011]。 图 2.2 提供了相应的 CUDA 版本的 SAXPY，它将执行拆分为 CPU 和 GPU。

图 2.2 中的示例演示了 CUDA 和相关编程模型（例如 OpenCL [Kaeli et al., 2015]）提供的抽象。 代码从函数 main() 开始执行。 为了让示例专注于 GPU 上计算的特定细节，我们省略了分配和初始化数组 x 和 y 的细节。 接下来，调用函数 saxpy_serial。 此函数将参数 n 中向量 x 和 y 中的元素数、参数 a 中的标量值以及用于表示向量 x 和 y 的数组指针作为输入参数。 该函数迭代数组 x 和 y 的每个元素。 在每次迭代中，第 4 行的代码使用循环变量 i 读取值 x[i] 和 y[i]，将 x[i] 乘以 a，然后加上 y[i]，然后用结果更新 x[i] . 为简单起见，我们省略了 CPU 如何使用函数调用结果的细节。

接下来，我们考虑 SAXPY 的 CUDA 版本。 与传统的 C 或 CCC 程序类似，图 2.2 中的代码通过在 CPU 上运行函数 main() 开始执行。 我们将首先突出显示特定于 GPU 执行的方面，而不是逐行浏览此代码。

在 GPU 上执行的线程是函数指定的计算内核的一部分。 在 CUDA 版本的 SAXPY 中，如图 2.2 所示，第 1 行的 CUDA 关键字 global 表示核函数 saxpy 将在 GPU 上运行。 在图 2.2 的示例中，我们并行化了图 2.1 中的“for”循环。 具体来说，第 4 行“for”循环的每次迭代图 2.1 中的原始 CPU 专用 C 代码被翻译成一个单独的线程，运行图 2.2 中第 3-5 行的代码。

一个计算内核通常由数千个线程组成，每个线程都从运行相同的函数开始。 在我们的示例中，CPU 在第 17 行使用 CUDA 的内核配置语法在 GPU 上开始计算。 内核配置语法看起来很像 C 中的函数调用，其中包含一些附加信息，用于指定包含在三尖括号 (<<<>>>) 之间的线程数。 构成计算内核的线程被组织成一个层次结构，该层次结构由由经线组成的线程块网格组成。 在 CUDA 编程模型中，单个线程执行操作数为标量值（例如，32 位浮点）的指令。 为了提高效率，典型的 GPU 硬件以锁步方式一起执行多组线程。 这些组被 NVIDIA 称为扭曲，被 AMD 称为波前。 NVIDIA 扭曲由 32 个线程组成，而 AMD 波前由 64 个线程组成。 经纱是 NVIDIA 将其分组为一个更大的单元，称为协作线程阵列 (CTA) 或线程块。第 17 行指示计算内核应该启动由 nblocks 线程块组成的单个网格，其中每个线程块包含 256 个线程。 CPU 代码传递给内核配置语句的参数被分发到 GPU 上正在运行的线程的每个实例。

当今许多移动设备片上系统将 CPU 和 GPU 集成到单个芯片中，就像当今笔记本电脑和台式计算机上的处理器一样。 然而，传统上，GPU 有自己的 DRAM 内存，今天对于用于机器学习的数据中心内的 GPU 而言，这种情况仍在继续。 我们注意到 NVIDIA 引入了统一内存，它从 CPU 内存透明地更新 GPU 内存，从 GPU 内存透明地更新 CPU 内存。 在启用统一内存的系统中，运行时和硬件负责代表程序员执行复制。 鉴于人们对机器学习的兴趣日益浓厚，并且本书的目标是了解硬件，因此在我们的示例中，我们考虑由程序员管理的独立 GPU 和 CPU 存储器的一般情况。

遵循许多 NVIDIA CUDA 示例中使用的样式，我们使用前缀 h_ 来命名分配在 CPU 内存中的内存的指针变量，并使用前缀 d_ 来命名分配在 GPU 内存中的内存的指针。 在第 13 行，CPU 调用 CUDA 库函数 cudaMalloc。 此函数调用 GPU 驱动程序并要求它在 GPU 上分配内存以供程序使用。 对 cudaMalloc 的调用将 d_x 设置为指向 GPU 内存区域，该区域包含足够的空间来保存 n 个 32 位浮点值。 在第 15 行，CPU 调用 CUDA 库函数 cudaMemcpy。 该函数调用 GPU 驱动程序，并要求它将 h_x 指向的 CPU 内存中的数组内容复制到 d_x 指向的 GPU 内存中的数组中。

最后让我们关注 GPU 上线程的执行。并行编程中采用的常见策略是为每个线程分配一部分数据。为了促进这种策略，GPU 上的每个线程都可以在线程块网格中查找自己的身份。在 CUDA 中执行此操作的机制使用网格、块和线程标识符。在 CUDA 中，网格和线程块具有 x、y 和 z 维度。当它执行时，每个线程在网格和线程块内都有一个固定的、唯一的非负整数 x、y 和 z 坐标组合。每个线程块在网格内都有 x、y 和 z 坐标。类似地，每个线程在线程块内都有 x、y 和 z 坐标。这些坐标的范围由内核配置语法（第 17 行）设置。在我们的示例中，没有指定 y 和 z 维度，因此所有线程的 y 和 z 线程块和线程坐标的值都为零。在第 3 行，threadIdx.x 的值标识了线程在其线程块内的 x 坐标，blockIdx.x 指示了线程块在其网格内的 x 坐标。值 blockDim.x 表示 x 维度中的最大线程数。在我们的示例中，blockDim.x 将评估为256，因为这是在第 17 行指定的值。表达式 blockIdx.x*blockDim.x + threadIdx.x 用于计算偏移量 i，以便在访问数组 x 和 y 时使用。正如我们将看到的，使用索引 i 我们为每个线程分配了一个唯一的 x 和 y 元素。

在很大程度上，编译器和硬件的结合使程序员能够忽略经线中线程执行的锁步性质。 编译器和硬件允许在独立执行的 warp 中出现每个线程。 在图 2.2 的第 4 行，我们将索引 i 的值与数组 x 和 y 的大小 n 进行比较。 i 小于 n 的线程执行第 5 行。图 2.2 中的第 5 行执行图 2.1 中原始循环的一次迭代。 在网格中的所有线程都完成后，计算内核在第 17 行之后将控制权返回给 CPU。在第 18 行，CPU 调用 GPU 驱动程序将 d_y 指向的数组从 GPU 内存复制回 CPU 内存。

SAXPY 示例中未说明的 CUDA 编程模型的一些其他细节，但我们将在稍后讨论，如下所示。

CTA 中的线程可以通过每个计算核心的暂存器内存有效地相互通信。这个 scrathpad 被 NVIDIA 称为共享内存。每个流式多处理器 (SM) 都包含一个共享内存。共享内存中的空间在该 SM 上运行的所有 CTA 之间分配。 AMD 的下一代图形核心 (GCN) 架构 [AMD, 2012] 包括一个类似的暂存存储器，AMD 将其称为本地数据存储 (LDS)。这些暂存存储器很小，每个 SM 的大小为 16-64 KB，并且作为不同的存储空间暴露给程序员。程序员使用源代码中的特殊关键字（例如，CUDA 中的“shared”）将内存分配到暂存器内存中。暂存器存储器充当软件控制的缓存。虽然 GPU 还包含硬件管理的缓存，但通过此类缓存访问数据可能会导致频繁的缓存未命中。当程序员可以以可预测的方式识别频繁重复使用的数据时，应用程序将从使用暂存存储器中受益。与 NVIDIA 的 GPU 不同，AMD 的 GCN GPU 还包括 GPU 上所有内核共享的全局数据存储 (GDS) 暂存内存。 Scratchpad 内存在图形应用程序中用于在不同的图形着色器之间传递结果。例如，LDS 用于在 GCN [AMD, 2012] 中在顶点和像素着色器之间传递参数值。

CTA 中的线程可以使用硬件支持的屏障指令有效地同步。 不同 CTA 中的线程可以通信，但通过所有线程都可以访问的全局地址空间进行通信。 在时间和能源方面，访问这个全局地址空间通常比访问共享内存更昂贵。

NVIDIA 在 Kepler 一代 GPU 中引入了 CUDA 动态并行 (CDP) [NVIDIA Corporation, a]。 CDP 的动机是观察到数据密集型不规则应用程序可能导致 GPU 上运行的线程之间的负载不平衡，从而导致 GPU 硬件未得到充分利用。 在许多方面，其动机类似于动态扭曲形成 (DWF) [Fung et al., 2007] 和第 3.4 节中讨论的相关方法。

2.2 GPU INSTRUCTION SET ARCHITECTURES

在本节中，我们将简要讨论计算内核从高级语言（例如 CUDA 和 OpenCL）到 GPU 硬件执行的汇编级别的转换以及当前 GPU 指令集的形式。 GPU 架构与 CPU 架构有些不同的一个有趣方面是 GPU 生态系统为支持指令集进化而发展的方式。 例如，x86 微处理器向后兼容 1976 年发布的 Intel 8086。向后兼容性意味着为上一代架构编译的程序将在下一代架构上运行而无需任何更改。 因此，40 年前为 Intel 8086 编译的软件理论上可以在当今的任何 x86 处理器上运行。

2.2.1 NVIDIA GPU INSTRUCTION SET ARCHITECTURES

鉴于有时有大量供应商提供 GPU 硬件（每个供应商都有自己的硬件设计），通过 OpenGL 着色语言 (OGSL) 和微软的高级着色语言 (HLSL) 的指令集虚拟化级别变得普遍，因为早期的 GPU 变得可编程。当 NVIDIA 在 2007 年初推出 CUDA 时，他们决定遵循类似的路径，并推出自己的用于 GPU 计算的高级虚拟指令集架构，称为并行线程执行 ISA，或 PTX [NVI，2017]。 NVIDIA 在 CUDA 的每个版本中都完整记录了这种虚拟指令集架构，以至于本书作者很容易开发出支持 PTX 的 GPGPU-Sim 模拟器 [Bakhoda et al., 2009]。 PTX 在许多方面类似于标准精简指令集计算机 (RISC) 指令集架构，如 ARM、MIPS、SPARC 或 ALPHA。它还与优化编译器中使用的中间表示有相似之处。一个这样的例子是使用一组无限的虚拟寄存器。图 2.3 说明了图 2.2 中 SAXPY 程序的 PTX 版本。

在 GPU 上运行 PTX 代码之前，有必要将 PTX 编译为硬件支持的实际指令集架构。 NVIDIA 将此级别称为 SASS，它是“Streaming ASSembler”的缩写 [Cabral, 2016]。从 PTX 到 SASS 的转换过程可以通过 GPU 驱动程序或提供的名为 ptxas 的独立程序来完成
使用 NVIDIA 的 CUDA 工具包。 NVIDIA 并未完整记录 SASS。虽然这使得学术研究人员更难开发捕获所有编译器优化效果的架构模拟器，但它使 NVIDIA 从客户需求中解放出来，在硬件级别提供向后兼容性，从而能够从一代到下一代完全重新设计指令集架构。不可避免地，希望在低层次了解性能的开发人员开始创建自己的工具来反汇编 SASS。第一次这样的努力归功于 Wladimir Jasper van der Laan 并命名为“decuda”[van der Lann]，于 2007 年底出现在 NVIDIA 的 GeForce 8 系列 (G80) 上，在第一次发布支持 CUDA 的硬件的大约一年内。 decuda 项目对 SASS 指令集有了足够详细的理解，因此可以开发汇编程序。这有助于在 GPGPU-Sim 3.2.2 [Tor M. Aamodt 等人] 中开发对 NVIDIA GT200 架构的 SASS 支持。 NVIDIA 最终推出了一个名为 cuobjdump 的工具，并开始部分记录 SASS。 NVIDIA 的 SASS 文档 [NVIDIA Corporation, c] 当前（2018 年 4 月）仅提供了汇编操作码名称的列表，但没有提供有关操作数格式或 SASS 指令语义的详细信息。最近，随着 GPU 用于机器学习的爆炸性增长以及对性能优化代码的需求，其他人已经为后续架构开发了类似于 decuda 的工具，例如 NVIDIA 的 Fermi [Yunqing] 和 NVIDIA 的 Maxwell 架构 [Gray]

图 2.4 说明了我们的 SAXPY 内核的 SASS 代码，该代码为 NVIDIA 的 Fermi 架构 [NVI，2009] 编译并使用 NVIDIA 的 cuobjdump（CUDA 工具包的一部分）提取。图 2.4 中的第一列是指令的地址。第二列是汇编，第三列是编码指令。如上所述，NVIDIA 仅部分记录了他们的硬件组装。比较图 2.3 和图 2.4，可以注意到虚拟和硬件 ISA 级别之间的异同。在高层次上有重要的相似之处，例如都是 RISC（都使用加载和存储来访问内存）和都使用预测 [Allen et al., 1983]。更细微的区别包括： (1) PTX 版本有一组基本上无限可用的寄存器，因此每个定义通常使用一个新寄存器，很像静态单分配 [Cytron et al., 1991] 而 SASS 使用一组有限的寄存器； (2) 内核参数通过可被 SASS 中的非加载/存储指令访问的存储常量内存传递，而参数则分配到 PTX 中它们自己单独的“参数”地址空间。

图 2.5 说明了由相同版本的 CUDA 生成的 SAXPY 的 SASS 代码，但用于 NVIDIA 的 Pascal 架构，并使用 NVIDIA 的 cuobjdump 提取。将图 2.5 与图 2.4 进行比较，很明显 NVIDIA 的 ISA 发生了重大变化，包括在指令编码方面。图 2.5 包含一些没有反汇编指令的行（例如，在第 3 行的地址 0x0000）。这些是 NVIDIA Kepler 架构中引入的特殊“控制指令”，以消除使用记分板进行显式依赖性检查的需要 [NVIDIA Corporation, b]。 Lai 和 Seznec [2013] 探索了 Kepler 架构的控制指令编码。正如 Lai 和 Seznec [2013] 所指出的，这些控制指令似乎类似于 Tera 计算机系统 [Alverson 等人，1990] 上的显式依赖前瞻。 Gray 描述了他们能够为 NVIDIA 的 Maxwell 架构推断出的控制指令编码的大量细节。根据 Gray 的说法，Maxwell 中每三个常规指令对应一个控制指令。如图 2.5 所示，NVIDIA 的 Pascal 架构似乎也是这种情况。根据 Gray 的说法，Maxwell 上的 64 位控制指令包含三组 21 位，对以下三个指令中的每一个的以下信息进行编码：停顿计数；产量提示标志；和写、读和等待依赖障碍。 Gray 还描述了在常规指令上使用寄存器重用标志，这也可以在图 2.5 中看到（例如，R0.reuse 用于第 7 行整数短乘加指令 XMAD 中的第一个源操作数）。这似乎表明从 Maxwell 开始在 NVIDIA GPU 中添加了“操作数重用缓存”（参见第 3.6.1 节中的相关研究）。这种操作数重用缓存似乎使每次主寄存器文件访问都可以多次读取寄存器值，从而降低能耗和/或改进的性能。

2.2.2 AMD GRAPHICS CORE NEXT INSTRUCTION SET ARCHITECTURE

与 NVIDIA 相比，当 AMD 推出其 Southern Islands 架构时，他们发布了完整的硬件级 ISA 规范 [AMD, 2012]。 Southern Islands 是第一代 AMD 下一代图形核心 (GCN) 架构。 AMD 硬件 ISA 文档的可用性帮助学术研究人员开发在较低级别工作的模拟器 [Ubal 等人，2012]。 AMD 的编译流程还包括一个称为 HSAIL 的虚拟指令集架构，作为异构系统架构 (HSA) 的一部分。

AMD 的 GCN 架构和 NVIDIA GPU（包括 NVIDIA 最新的 Volta 架构 [NVIDIA Corp., 2017]）之间的一个主要区别是单独的标量和向量指令。图 2.6 和 2.7 再现了来自 AMD [2012] 的高级 OpenCL（类似于 CUDA）代码和 AMD 南岛架构的等效机器指令的示例。在图 2.7 中，标量指令以 s_ 开头，向量指令以 v_ 开头。在 AMD GCN 架构中，每个计算单元（例如 SIMT 核心）包含一个标量单元和四个向量单元。矢量指令在矢量单元上执行，并为波前的每个单独线程计算不同的 32 位值。相比之下，在标量单元上执行的标量指令计算波前中所有线程共享的单个 32 位值。在图 2.7 所示的示例中，标量指令与控制流处理相关。特别是，exec 是一个特殊寄存器，用于为 SIMT 执行预测单个向量通道的执行。第 3.1.1 节更详细地描述了在 GPU 上使用掩码进行控制流处理。 GCN 架构中标量单元的另一个潜在好处是，SIMT 程序中计算的某些部分经常会计算出与线程 ID 无关的相同结果（参见第 3.5 节）。

AMD 的 GCN 硬件指令集手册 [AMD, 2012] 提供了许多关于 AMD GPU 硬件的有趣见解。 例如，为了为长延迟操作启用数据依赖解析，AMD 的 GCN 架构包括 S_WAITCNT 指令。 对于每个波前，有三个计数器：矢量存储器计数、本地/全局数据存储计数和寄存器导出计数。 这些中的每一个都表示给定类型的未完成操作的数量。 编译器或程序员插入 S_WAITCNT 指令以使波前等待，直到未完成操作的数量减少到指定阈值以下。