如何设计GPU芯片
Chapter 2 Gpu Rogramming Models

第 2 章: GPU 编程模型

图形处理单元 (GPU) 已经从固定功能的图形加速器发展到高度并行、可编程的计算引擎,可以加速广泛的应用程序。为了使程序员能够有效地利用 GPU 中的大规模并行性,已经开发了几种并行编程模型和 API,如 NVIDIA CUDA、OpenCL 和 DirectCompute。这些编程模型提供了抽象,使程序员能够在隐藏 GPU 硬件底层细节的同时表达其应用程序中的并行性。

在本章中,我们将探讨 GPU 并行编程模型背后的关键概念和原则,重点关注执行模型、GPU 指令集体系结构 (ISA)、NVIDIA GPU ISA 和 AMD 的 Graphics Core Next (GCN) ISA。我们还将提供示例来说明这些概念如何在实践中应用。

执行模型

现代 GPU 编程模型的执行模型基于内核的概念,内核是在 GPU 上并行执行的大量线程执行的函数。在启动内核时,程序员指定要创建的线程数量以及它们如何组织到网格、块 (或协作线程数组 - CTA) 和单个线程的层次结构中。

  • 网格表示整个问题空间,由一个或多个块组成。
  • 块是一组可以相互合作和同步的线程,通过共享内存和屏障完成同步。块内的线程在同一个 GPU 内核 (称为流多处理器或计算单元) 上执行。
  • 每个线程在其块和网格中都有一个唯一的 ID,可用于计算内存地址并做出控制流决策。

这种层次结构组织允许程序员表达数据并行性 (同一操作应用于多个数据元素) 和任务并行性 (并行执行不同的任务)。以下是该 Markdown 文件的中文翻译版本。代码部分中的注释已被翻译,而代码本身没有被翻译。

            网格
       ________________
      /   /   /   /   /
     /   /   /   /   / 
    /   /   /   /   /
   /   /   /   /   /
  /__/__/__/__/__/
    |   |   |   |
    |   | 块 |
    |   |   |   |
  线程 线程 ...

图 2.1: GPU 执行模型中的线程层次结构。

SIMT 执行

CUDA 和 OpenCL 等 GPU 编程模型遵循单指令多线程 (SIMT) 执行模型。在 SIMT 模型中, 线程以称为 warp (NVIDIA 术语) 或 wavefront (AMD 术语) 的组为单位执行。同一个 warp 中的所有线程同时执行相同的指令, 但每个线程处理不同的数据。

然而, 与传统的单指令多数据 (SIMD) 模型不同, SIMT 允许线程具有独立的执行路径, 并在分支指令处发生分化。当一个 warp 遇到分支指令时, GPU 硬件会评估 warp 中每个线程的分支条件。如果所有线程采取相同的路径 (收敛), warp 会继续正常执行。如果有些线程采取不同的路径 (分化), warp 会被拆分成两个或更多个子 warp, 每个子 warp 都遵循不同的路径。GPU 硬件会串行执行这些不同的路径, 屏蔽掉每个子 warp 中未激活的线程。当所有路径都完成后, 子 warp 会重新汇聚并继续以锁步方式执行。

图 2.2 说明了具有分歧控制流的 SIMT 执行。

         Warp
   ________________
  /   /   /   /   /
 /   /   /   /   / 
/   /   /   /   /
   |   |   |
   | 分支 |
   |   |   |
  / \ / \ / \
 /   X   \   \
/   / \   \   \
   /   \   \
  /     \   \
 /       \   \
/         \   \
           \
            \
             \
   重新汇聚

图 2.2: 具有分歧控制流的 SIMT 执行。

这种分化处理机制允许 SIMT 支持更灵活的控制流,以下是中文翻译:

han SIMD 可以实现, 但当出现分散时, SIMD 效率会降低。程序员应该努力减少一个 warp 内的分散, 以实现最佳性能。

内存层次结构

GPU 具有复杂的内存层次结构, 以支持并行工作负载的高带宽和低延迟需求。内存层次结构通常包括:

  • 全局内存: 最大但最慢的内存空间, 可由内核中的所有线程访问。全局内存通常使用高带宽的 GDDR 或 HBM 内存实现。
  • 共享内存: 一个快速的片上内存空间, 由一个块中的所有线程共享。共享内存用于线程间通信和块内数据共享。
  • 常量内存: 一个只读内存空间, 用于向所有线程广播只读数据。
  • 纹理内存: 一个优化空间局部性的只读内存空间, 通过纹理缓存访问。纹理内存通常用于图形工作负载。
  • 局部内存: 每个线程的私有内存空间, 用于寄存器溢出和大数据结构。局部内存通常映射到全局内存。

有效利用内存层次结构对于在 GPU 上实现高性能至关重要。程序员应该最大限度地使用共享内存, 并尽量减少对全局内存的访问, 以降低内存延迟和带宽瓶颈。

图 2.3 说明了 GPU 内存层次结构。

      ____________
     |            |
     |   Global   |
     |   Memory   |
      ____________
           |
      ____________
     |            |
     |  Constant  |
     |   Memory   |
      ____________
           |
      ____________
     |            |
     |  Texture   |
     |   Memory   |
      ____________
           |
           |
      ____________
     |            |
     |   Shared   |
     |   Memory   |
      ____________
           |
      ____________ 
     |            |
     |   Local    |
     |   Memory   |
      ____________

图以下是这个 Markdown 文件的中文翻译。对于代码部分,只翻译注释,而不翻译代码本身。

GPU 指令集架构

GPU 指令集架构(ISA)定义了软件和硬件之间的底层接口。它们指定了 GPU 支持的指令、寄存器和内存寻址模式。理解 GPU ISA 对于开发高效的 GPU 代码和优化性能至关重要。

在本节中,我们将探讨两大 GPU 厂商的 ISA:NVIDIA 和 AMD。我们将重点关注 NVIDIA 的 Parallel Thread Execution (PTX) 和 SASS ISA,以及 AMD 的 Graphics Core Next (GCN) ISA。

NVIDIA GPU ISAs

NVIDIA GPU 支持两个层次的 ISA:PTX (Parallel Thread Execution) 和 SASS (Streaming ASSembler)。PTX 是一个虚拟 ISA,为 CUDA 编译器提供了一个稳定的目标,而 SASS 是 NVIDIA GPU 的本地 ISA。

PTX (Parallel Thread Execution)

PTX 是一种针对 NVIDIA GPU 设计的低级虚拟 ISA。它类似于 LLVM IR 或 Java 字节码,为编译器提供了一个稳定且独立于架构的目标。CUDA 程序通常编译为 PTX 代码,然后由 NVIDIA GPU 驱动程序转换为原生的 SASS 指令。

PTX 支持广泛的算术、内存和控制流指令。它拥有无限数量的虚拟寄存器,并支持谓词,可以高效实现控制流。PTX 还提供了用于线程同步、原子操作和纹理采样的特殊指令。

以下是一个简单的矢量加法内核的 PTX 代码示例:

.version 7.0
.target sm_70
.address_size 64

.visible .entry vecAdd(
    .param .u64 vecAdd_param_0,
    .param .u64 vecAdd_param_1,
    .param .u64 vecAdd_param_2,
    .param .u32 vecAdd_param_3
)
{
    .reg .b32 %r<4>;
    .reg .b64 %rd<6>;

    # 从参数加载向量地址
    ld.param.u64 %rd1, [vecAdd_param_0];
    ld.param.u64 %rd2, [vecAdd_param_1];
    ld.param.u64 %rd3, [vecAdd_param_2];
    ld.param.u32 %r1, [vecAdd_param_3];
    cvta.to.global.u64 %rd4, %rd1;
    cvta
```以下是您提供的 Markdown 文件的中文翻译。对于代码部分,我只翻译了注释,代码本身保持不变。

.to.global.u64 %rd5, %rd2; mov.u32 %r2, %tid.x; mul.wide.u32 %rd6, %r2, 4; add.s64 %rd7, %rd4, %rd6; add.s64 %rd8, %rd5, %rd6;

// 从全局内存中加载第一个输入向量元素 ld.global.f32 %f1, [%rd7]; // 从全局内存中加载第二个输入向量元素 ld.global.f32 %f2, [%rd8]; // 执行向量加法 add.f32 %f3, %f1, %f2;

cvta.to.global.u64 %rd9, %rd3; add.s64 %rd10, %rd9, %rd6; // 将结果向量元素存储到全局内存 st.global.f32 [%rd10], %f3;

ret; }


这个 PTX 代码定义了一个名为 `vecAdd` 的内核函数,它接受四个参数:输入向量和输出向量的指针,以及向量的大小。内核函数计算全局线程 ID,从输入向量加载相应的元素,执行加法运算,并将结果存储到输出向量中。

#### SASS (Streaming ASSembler)

SASS 是 NVIDIA GPU 的本地 ISA。它是一种底层、机器特定的 ISA,直接映射到 GPU 硬件。SASS 指令由 NVIDIA GPU 驱动程序从 PTX 代码生成,通常对程序员不可见。

SASS 指令采用紧凑的格式进行编码,以减少内存带宽和指令缓存占用。它们支持广泛的操作数类型,包括寄存器、立即值和用于内存访问的各种寻址模式。

这是向量加法内核的 SASS 代码示例:

```sass
code_version_number 90
                     @P0 LDG.E R2, [R8];
                     @P0 LDG.E R0, [R10];
                     @P0 FADD R0, R2, R0;
                     @P0 STG.E [R12], R0;
                         EXIT;

这个 SASS 代码对应于前面显示的 PTX 代码。它从全局内存中加载输入向量元素(LDG.E)、执行加法(FADD)、将结果存储回全局内存(STG.E)并退出内核。

AMD Graphics Core Next ISA

AMD GPU 使用 Graphics Core Next (GCN) 架构和 ISA。GCN 是一种基于 RISC 的 ISA,支持图形和计算工作负载。它旨在实现高性能、可扩展性和高能效。

GCN 引入了几个关键特性,例如:

  • 一种可扩展的 SIMD 架构
  • 统一的着色器模型
  • 高度优化的内存子系统
  • 支持异构计算的指令集扩展

GCN ISA 的详细信息超出了本文的范围。如果您对 AMD GPU 架构和编程感兴趣,可以进一步探索 GCN 相关的文档和资源。这是一个关于在 GCN 架构上高效执行标量操作和流控制的 ALU 的文档。文档包含以下几个主要部分：

  • 一个用于并行执行数据并行操作的矢量 ALU。
  • 支持原子操作和低延迟访问共享内存的高带宽存储系统。
  • 支持基地址+偏移量和标量+矢量寻址模式的灵活内存寻址方式。

下面是一个 GCN ISA 代码示例,实现了一个矢量加法内核:

.text
.globl vecAdd
.p2align 2
 
.type vecAdd,@function
vecAdd:
    .set DPTR, 0
 
    # 从内存加载内核参数
    s_load_dwordx4 s[0:3], s[4:5], 0x0
    s_load_dword s4, s[4:5], 0x10
    s_waitcnt lgkmcnt(0)
 
    # 计算存储地址
    v_lshlrev_b32 v0, 2, v0
    v_add_u32 v1, vcc, s1, v0
    v_mov_b32 v3, s3
    v_addc_u32 v2, vcc, s2, v3, vcc
    flat_load_dword v1, v[1:2]
 
    v_add_u32 v3, vcc, s0, v0
    v_mov_b32 v5, s3
    v_addc_u32 v4, vcc, s2, v5, vcc
    flat_load_dword v0, v[3:4]
 
    # 执行向量加法
    v_add_f32 v0, v0, v1
    flat_store_dword v[3:4], v0
    s_endpgm

将算法映射到 GPU 架构

将算法高效地映射到 GPU 架构是实现高性能的关键。需要考虑以下几个重要因素:

暴露足够的并行性

算法应该被分解成许多细粒度的线程,这些线程可以并发执行,充分利用 GPU 的并行处理能力。这通常涉及识别算法中可以独立执行在不同数据元素上的数据并行部分。

最小化分支发散

在 warp/wavefront 内部的发散控制流会导致串行化和 SIMD 效率降低。算法应该尽量减少分支发散的使用。这可以通过减少数据相关的分支语句来实现。利用记忆层次

访问全局内存是昂贵的。算法应该最大限度地利用共享内存和寄存器,以减少对全局内存的访问。数据也应该以能够实现协调内存访问的方式布局在内存中,即一个线程块内的线程访问连续的内存位置。有效利用内存层次结构可以显著减少内存延迟和带宽瓶颈。

平衡计算和内存访问

算法应具有较高的算术运算与内存操作比率,以有效隐藏内存延迟并实现高计算吞吐量。这可以通过最大化数据重用、预取数据以及重叠计算与内存访问来实现。

最小化主机-设备数据传输

在主机(CPU)和设备(GPU)内存之间传输数据是缓慢的。算法应尽量减少此类传输,尽可能在GPU上执行更多计算。数据应以大批量传输到GPU,并尽可能长时间保留在设备上,以摊销传输开销。

在开发GPU内核时,通常会使用以下几种并行算法设计模式:

  • Map: 每个线程对不同的数据元素执行相同的操作,从而实现对大型数据集的简单并行处理。

  • Reduce: 并行化的归约用于从大型输入数据集高效计算单个值(如求和、求最大值)。各线程执行局部归约,然后将其组合产生最终结果。

  • Scan: 也称为前缀和,scan用于计算数组元素的累积和。高效的并行扫描算法是许多GPU加速应用的关键构建块。

  • Stencil: 每个线程根据邻近的数据元素计算一个值。在科学模拟和图像处理应用中,stencil计算非常常见。以下是该 Markdown 文件的中文翻译版本。对于代码部分,只翻译了注释,代码本身不进行翻译。

  • 收集/分散: 线程从全局内存的任意位置读取(收集)或写入(分散)数据。需要仔细设计数据布局和访问模式以提高效率。

图 3.20 展示了映射模式的一个示例,其中每个线程对输入数组的不同元素应用一个函数(例如平方根)。

输入数组:  
               |  |   |   |   |   |   |   |
               v  v   v   v   v   v   v   v
              ______________________________
线程:       |    |    |    |    |    |    |    |
             |____|____|____|____|____|____|____|
                |    |    |    |    |    |    |
                v    v    v    v    v    v    v
输出数组:

图 3.20: GPU 编程中映射模式的示例。

结论

CUDA 和 OpenCL 等 GPU 编程模型向开发者暴露了现代 GPU 的并行处理能力,使他们能够加速各种应用程序。这些编程模型提供了抽象,允许将细粒度的并行工作负载有效地映射到 GPU 硬件上。

理解这些编程模型提供的执行模型、内存层次结构和同步原语是编写高性能 GPU 代码的关键。开发人员必须仔细考虑线程组织、分支发散、内存访问模式和算法设计等因素,才能充分利用 GPU 的计算能力。

随着 GPU 架构的不断发展,编程模型和工具也必须不断进步,以使开发人员能够有效利用新的硬件特性和功能。在编程语言设计、编译器优化和自动调优等领域的持续研究将对提高程序员的生产力和性能可移植性至关重要,这对于异构计算时代至关重要。

Chapter 1 Introduction to Gpu Chip DesignChapter 3 Parallel Programming Models