General-Purpose Graphics Processor Architecture机翻（二 ...

yukamu · 发表于 2022-5-19 18:10

本章的目标是提供足够的上下文来说明如何为非图形计算编程 GPU 编程，以便那些没有 GPU 经验的人可以关注后面章节的讨论。我们在这里专注于基本材料，将更深入的报道留给其他参考资料（例如，[Kirk and Wen-Mei, 2016]）。存在许多可用于架构研究的 GPU 计算基准套件。学习如何对 GPU 进行编程与对 GPU 计算感兴趣的计算机架构师有关，以便更好地了解硬件/软件接口，但如果您想探索对硬件/软件接口进行更改作为研究的一部分，则它变得至关重要。在后一种情况下，现有基准可能不存在，因此可能需要通过修改现有 GPU 计算应用程序的源代码来创建。例如，探索在 GPU 上引入事务内存 (TM) 的研究需要这样做，因为当前的 GPU 不支持 TM（参见第 5.3 节）。
现代 GPU 采用广泛的 SIMD 硬件来利用 GPU 应用程序中的数据级并行。 CUDA 和 OpenCL 等 GPU 计算 API 不是直接将 SIMD 硬件暴露给程序员，而是具有类似 MIMD 的编程模型，允许程序员在 GPU 上启动大量标量线程。这些标量线程中的每一个都可以遵循其独特的执行路径，并且可以访问任意内存位置。在运行时，GPU 硬件在 SIMD 硬件上执行称为warps（或 AMD 术语中的wavefronts）的标量线程组，以利用它们的规律性和空间局部性。这种执行模型称为单指令多线程 (SIMT) [Lindholm et al., 2008a, Nickolls and Reusch, 1993]。
本章的其余部分扩展了这一讨论，并组织如下。在第 2.1 节中，我们探讨了最近 GPU 编程模型使用的概念执行模型，并简要总结了过去十年发布的典型 GPU 的执行模型。在 2.2 节中，我们探讨了 GPU 计算应用程序的编译过程，并简要介绍了 GPU 指令集架构。
2.1 执行模型

GPU 计算应用程序开始在 CPU 上执行。对于分离式 GPU，应用程序的 CPU 部分通常会分配内存用于 GPU 上的计算，然后将输入数据传输到 GPU 内存，最后在 GPU 上启动计算内核。对于集成式 GPU，只需要最后一步。计算内核由（通常）数千个线程组成。每个线程执行相同的程序，但可能会根据计算结果遵循通过该程序的不同控制流。下面我们使用一个用 CUDA 编写的特定代码示例详细考虑这个流程。在下一节中，我们将在汇编级别查看执行模型。我们的讨论并没有停留在 GPU 编程模型的性能方面。然而，Seo 等人提出了一个有趣的观察结果。 [2011] 在 OpenCL（一种类似于 CUDA 的编程模型，可以编译到许多架构）的上下文中，为一种架构（例如 GPU）精心优化的代码可能在另一种架构（例如 CPU）上表现不佳。
图 2.1 为众所周知的单精度标量值 A 乘以向量值 X 加上向量值 Y（称为 SAXPY）的 CPU 实现提供了 C 代码。 SAXPY 是著名的基本线性代数软件 (BLAS) 库 [Lawson 等人，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 和 y，将 x 乘以 a，然后加上 y，然后用结果更新 x .为简单起见，我们省略了 CPU 如何使用函数调用结果的细节。

接下来，我们考虑 SAXPY 的 CUDA 版本。与传统的 C 或 C++ 程序类似，图 2.2 中的代码通过在 CPU 上运行函数 main() 开始执行。我们将首先强调 GPU 执行的特定方面，而不是逐行浏览此代码。
在 GPU 上执行的线程是由函数指定的计算内核的一部分。在 SAXPY 的 CUDA 版本中，如图 2.2 所示，第 1 行的 CUDA 关键字 __global__ 表示内核函数 saxpy 将在 GPU 上运行。在图 2.2 的示例中，我们并行化了图 2.1 中的“for”循环。具体来说，图 2.1 中的原始 CPU-only C 代码中第 4 行“for”循环的每次迭代都被转换为运行图 2.2 中第 3-5 行代码的单独线程。

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

译者注：以NVIDIA为例，插入一段上述概念的解释。下面的多张图均截取自 NVIDIA H100 Tensor Core GPU Architecture（NVIDIA H100 白皮书）。
译者注：图14左图展示了A100中，Grid由Thread Blocks组成，而Thread Block由Threads组成。图中没有画出Warp（线程束）的概念。实际上，Warp是硬件最小的执行粒度，WarpPerThreadBlock = ceil(ThreadPerThreadBlock / WarpSize)，其中WarpSize为32。
译者注：图14右图中Thread Block Cluster为最新的H100引入的新的层级，旧的GPU架构中没有这一层级。本书成书于H100之前，因此没有提及这个层级。

译者注：下图中Threads / Warp = 32，表示每个Warp包含32个Threads。

当今的许多移动设备片上系统都将 CPU 和 GPU 集成到单个芯片中，就像当今笔记本电脑和台式计算机上的处理器一样。然而，传统上，GPU 有自己的 DRAM 内存，今天在用于机器学习的数据中心内发现的 GPU 仍然如此。我们注意到 NVIDIA 引入了统一内存，它透明地从 CPU 内存更新 GPU 内存，以及从 GPU 内存更新 CPU 内存。在启用统一内存的系统中，运行时 (runtime) 和硬件负责代表程序员执行拷贝。鉴于对机器学习的兴趣日益增加，并且本书的目标是理解硬件，在我们的示例中，我们考虑由程序员管理的独立 GPU 和 CPU 内存的一般情况。
遵循许多 NVIDIA CUDA 示例中使用的风格，我们使用前缀 h_ (host) 来命名分配在 CPU 内存中的内存的指针变量，使用 d_ (device) 来命名分配在 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（被NVIDIA称为thread block）中的线程可以通过每个计算核心暂存器 (scratchpad memory) 内存有效地相互通信。这个 scratchpad 被 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] 中的顶点和像素着色器之间传递参数值。

译者注：仍然引用 NVIDIA H100 Tensor Core GPU Architecture（NVIDIA H100 白皮书），解释Streaming Multiprocessor的概念。
译者注：如图6所示，每个GPU (device) 包含多个SMs。

译者注：GH100 单个 SM 内部的细节如图7所示。我们可以看到，每个 SM 都包含一个 Shared Memory。注意 L1 Data Cache 和 Shared Memory 的区别：L1 Data Cache 由硬件管理，而 Shared Memory 由软件管理。

CTA 中的线程可以使用硬件支持的屏障指令有效地同步。不同 CTA 中的线程可以通信，但需要通过所有线程都可以访问的全局地址空间进行通信。就时间和精力而言，访问这个全局地址空间通常比访问共享内存更昂贵。
NVIDIA 在 Kepler 一代 GPU 中引入了 CUDA 动态并行 (CDP) [NVIDIA Corporation, a]。 CDP 的动机是观察到数据密集型不规则应用程序可能导致 GPU 上运行的线程之间的负载不平衡，从而导致 GPU 硬件未被充分利用。在许多方面，其动机类似于 Dynamic Warp Formation (DWF) [Fung et al., 2007] 以及第 3.4 节中讨论的相关方法。
2.2 GPU 指令集架构

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

一直以来，大量供应商提供 GPU 硬件（每个都有自己的硬件设计），通过 OpenGL 着色语言 (OGSL) 和微软的高级着色语言 (HLSL) ，一定水平的指令集虚拟化随着随着早期的 GPU 变得可编程而变得很常见。当 NVIDIA 在 2007 年初推出 CUDA 时，他们决定走类似的道路，并为 GPU 计算引入了自己的高级虚拟指令集架构，称为并行线程执行 (Parallel Thread Execution) 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 版本。

NVIDIA 的 PTX ISA：Parallel Thread Execution ISA (v7.0)

在 GPU 上运行 PTX 代码之前，有必要将 PTX 编译为硬件支持的实际指令集架构。 NVIDIA 将此级别称为 SASS，它是“Streaming ASSembler”的缩写 [Cabral，2016]。从 PTX 转换为 SASS 的过程可以通过 GPU 驱动程序或 NVIDIA 的 CUDA 工具包提供的名为 ptxas 的独立程序来完成。 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 等人] 中开发对 SASS 的支持，直至 NVIDIA 的 GT200 架构。 NVIDIA 最终推出了一个名为 cuobjdump 的工具，并开始部分记录 SASS。 NVIDIA 的 SASS 文档 [NVIDIA Corporation, c] 当前（2018 年 4 月）仅提供了汇编操作码名称的列表，但没有提供有关操作数格式或 SASS 指令语义的详细信息。最近，随着在机器学习中使用 GPU 的爆炸式增长以及对性能优化代码的需求，其他人已经为后续架构开发了类似于 decuda 的工具，例如 NVIDIA 的 Fermi [Yunqing] 和 NVIDIA 的 Maxwell 架构 [Gray]。
图 2.4 展示了我们为 NVIDIA 的 Fermi 架构 [NVI, 2009] 编译并使用 NVIDIA 的 cuobjdump（CUDA 工具包的一部分）提取的 SAXPY 内核的 SASS 代码。图 2.4 中的第一列是指令的地址。第二列是汇编，第三列是编码指令。如上所述，NVIDIA 仅部分记录了他们的硬件组装。比较图 2.3 和图 2.4，可以看出虚拟和硬件 ISA 级别之间的相似之处和不同之处。在高层次上存在重要的相似之处，例如都是 RISC（都使用加载和存储来访问内存）和都使用谓词 [Allen et al., 1983]。更细微的区别包括： (1) PTX 版本具有基本上无限的可用寄存器集，因此每个定义通常使用一个新寄存器，很像静态单一分配 [Cytron 等人，1991]，而 SASS 使用有限的一组寄存器； (2) 内核参数通过存储在 SASS 中的非加载/存储指令可以访问的常量内存传递，而参数在 PTX 中分配到它们自己单独的“参数”地址空间中。

图 2.5 展示了 SAXPY 的 SASS 代码，该代码由相同版本的 CUDA 但针对 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 et al., 1990]。 Gray 描述了他们能够为 NVIDIA 的 Maxwell 架构推断出的控制指令编码的大量细节。根据 Gray 的说法，Maxwell 中每三个常规指令就有一个控制指令。这似乎也适用于 NVIDIA 的 Pascal 架构，如图 2.5 所示。根据 Gray 的说法，Maxwell 上的 64 位控制指令包含三组 21 位，为以下三个指令中的每一个编码以下信息：停顿计数；产量提示标志；以及写入、读取和等待依赖屏障。 Gray 还描述了寄存器重用标志 (register reuse flags) 在常规指令上的使用，如图 2.5 所示（例如，R0.reuse 用于第 8 行的整数短乘加指令 (Integer Short Multiply Add instruction) XMAD 中的第一个源操作数）。这似乎表明从 Maxwell 开始在 NVIDIA GPU 中添加了“操作数重用缓存 (operand reuse cache)”（参见第 3.6.1 节中的相关研究）。这种操作数重用缓存似乎能够为每个主寄存器文件访问多次读取寄存器值，从而降低能耗和/或提高性能。

2.2.2 AMD 显卡核心下一代指令集架构

与 NVIDIA 相比，AMD 推出了他们的 Southern Islands 架构，他们发布了完整的硬件级 ISA 规范 [AMD, 2012]。 Southern Islands 是 AMD 的第一代 Graphics Core Next (GCN) 架构。 AMD 硬件 ISA 文档的可用性帮助学术研究人员开发了在较低级别上工作的模拟器 [Ubal et al., 2012]。 AMD 的编译流程还包括一个称为 HSAIL 的虚拟指令集架构，作为异构系统架构 (HSA) 的一部分。
AMD 的 GCN 架构和 NVIDIA GPU（包括 NVIDIA 最新的 Volta 架构 [NVIDIA Corp., 2017]）之间的一个关键区别是单独的标量和向量指令。图 2.6 和 2.7 再现了 AMD [2012] 的高级 OpenCL（类似于 CUDA）代码示例和 AMD Southern Islands 架构的等效机器指令。在图 2.7 中，标量指令以 s_ 开头，向量指令以 v_ 开头。在 AMD GCN 架构中，每个计算单元（例如 SIMT 核心）都包含一个标量单元和四个向量单元。向量指令在向量单元上执行，并为 wavefront 中的每个单独线程计算不同的 32 位值。相反，在标量单元上执行的标量指令计算 wavefront 中所有线程共享的单个 32 位值。在图 2.7 所示的示例中，标量指令与控制流处理有关。特别是， exec 是一个特殊寄存器，用于预测 SIMT 执行的各个向量通道的执行。在第 3.1.1 节中更详细地描述了在 GPU 上使用掩码 (masking) 进行控制流处理。 GCN 架构中标量单元的另一个潜在好处是，SIMT 程序中计算的某些部分经常会计算相同的结果，而与线程 ID 无关（参见第 3.5 节）。

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

AMD 的 GCN 硬件指令集手册：Southern Islands Series Instruction Set Architecture

		自动登录	找回密码
密码			立即注册

General-Purpose Graphics Processor Architecture机翻（二 ...

本帖子中包含更多资源