GPU Architecture & CUDA Programming

一文带你厘清 GPU 的前世今生

这节课非常非常重要且硬核。笔者在9月中旬花1h30min听了一遍课程，完全没听懂。

在10月中旬重新仔细听了一遍，花了接近4h，现在应该说对这节课完全理解吃透了。

因此本章的整理内容会偏细致，复盘过程一共耗费了笔者约12h，大有所获！本文力求逻辑的连贯性。

Outline

1. 历史: 图形化处理器, 在最初只是为了加速3D游戏/图形渲染, 是如何一步一步走到今天的位置的: 大型应用的高度并行化计算引擎
   1. deep learning
   2. computer vision
   3. scientific computing
2. 如何使用 CUDA Language 对GPU进行编程
3. GPU架构解析

Graphics 101 + "Ancient" GPU

概念普及:

1. 图形渲染任务 (rendering task): computing how each triangle in 3D mesh contributes to appearance of each pixel in the image?
2. GPU最初设计意图: 加速 "实时渲染高复杂度的 2D/3D 场景"
   • 渲染: 是由 图形渲染工具(OpenGL) 完成的
   • 加速: 是由 GPU 完成的

历史发展:

(1) 图形如何渲染

1. 将建模物品用 2D/3D 的 triangle mesh 进行表示
2. 图形渲染工具(eg. OpenGL)起关键作用:
   • 给定一个triangle mesh, 它该放在哪个位置
   • 给每个triangle mesh里的 pixel, 计算并填充颜色
3. 图形渲染工具的可编程化:
   • 使用 GLSL (OpenGL shading language) shader program 来定义: 对每个pixel的处理行为

(2) 为什么需要GPU

GPU的特色是: Multi-core + SIMD + Multi-thread

因此, 它可以 为上述 shader program 提供 efficient execution

tldr: GPUs are very fast processors for performing the same computation (shader programs) in parallel on large collections of data (streams of vertices, fragments, and pixels)

因此, 可以认为GPU是一个 data-parallel programming system

(3) GPGPU [2002-2003]

GPGPU = “general purpose” computation on GPUs

由于当时GPU的唯一硬件接口是图形管线，如果想要利用GPU进行通用计算（GPGPU），程序员必须采用一种“Hack”方式:

1. 设定两个triangle mesh, 用来覆盖图像
2. 这样, 每个像素的着色器计算就相当于对输出数组中的一个元素进行计算

此时GPU可以加速, 除了图形渲染以外的其他功能, 如: 稀疏矩阵求解, 光线追踪...

(4) Brook stream programming language [2004]

Brook编译器充当了 高级通用计算模型 与 低级图形硬件接口 之间的桥梁

• 输入: 抽象的通用流程序，例如scale 内核（kernel）
  • 这种内核不是由OpenGL提供的，而是程序员使用Brook语言编写的，用于表达数据并行操作（例如，对所有流元素进行缩放）
• Brook编译器的转换: Brook编译器将这个通用流程序翻译成图形命令（例如 drawTriangles）和一套可以在当时的GPU上运行的图形着色器程序

[1] Brook编译器有啥用

Brook编译器的工作是必要的，因为在2007年之前，GPU硬件只能执行图形管线计算

通过将通用计算（如向量缩放）hack成图形渲染任务（如绘制覆盖屏幕的三角形，并使用片段着色器执行缩放），Brook才能够让GPU运行非图形程序

[2] Brook编译器, OpenGL, GPU 什么关系? Workflow/Dataflow是什么

OpenGL是一个API，用于将图形命令和着色器程序发送给GPU

Brook编译器生成的“图形命令”（如 drawTriangles）和“图形着色器程序 (aka. graphics shader programs)” 实际上就是 通过OpenGL（或类似的图形API）来驱动GPU执行 的底层指令

Brook抽象了底层复杂的图形编程步骤，让程序员能够专注于数据并行问题，而不是如何“欺骗”图形管线来执行计算

tldr: 高级程序员意图 --> Brook编译器 --> 图形指令+图形着色器程序 --> OpenGL API --> GPU

GPU Computer Mode

New: Computer-Mode API

(1) 在2007年之前, GPU专用于图形渲染, 其他用途(GPGPU)则需要hack

在这个时代, GPU 是专门用来做图形渲染加速的 —> GPU硬件 与 上层软件 之间的唯一接口, 就是 图形应用API (如 OpenGL API)

这也呼应了上面我们提到: GPGPU 需要 ”hack” 的原因

(2) 2007年, Nvidia Tesla 横空出世, GPU的通用计算架构登上历史舞台

在 2007 年之前，访问 GPU 硬件的唯一接口是图形管线（graphics pipeline）。如果用户想要在 GPU 的可编程核心上运行非图形程序（即通用计算，如 GPGPU 2002-2003 年代），他们必须使用像 Brook 这样的语言，将通用流程序翻译成图形命令（如 drawTriangles）和图形着色器程序，以“欺骗”图形管线执行计算

CUDA 和 Tesla 架构的引入 提供了第一个替代性的硬件接口 ，即 计算模式 (compute-mode) 接口

[1] 非图形化应用程序 现在可以更好地与 GPU 互动了:

1. 可以对 GPU 的 memory/buffer 进行操作
2. 可以通过 Graphics Driver 为GPU提供 a single Kernel Program Binary
3. 可以告诉 GPU 以 SPMD 的方式运行 kernel

[2] 在 2007 年后，GPU 同时拥有两条主要的接口API:

1. 图形路径 (OpenGL/DirectX): 用于 3D 渲染, 通过图形驱动程序发送着色器程序和 draw 命令, 利用 GPU 的并行核心来高效执行着色器程序
   • "原图形 pipeline 接口保留"
2. 计算路径 (CUDA/OpenCL): 用于通用并行计算, 通过计算模式接口发送内核程序和 launch 命令, 直接利用 GPU 的可编程核心
   • "新引入 computer-mode 接口"

OpenGL 与 GPU 的关系 (图形路径)

OpenGL 仍然是驱动 GPU 执行图形渲染任务的关键 API

• 图形管线仍然存在: 图形管线接口仍然是驱动 GPU 执行的路径之一
• 非可编程功能: GPU 上的许多有趣的非可编程功能仍然存在，用于加速图形管线操作的执行
• 并行运行: 当运行 CUDA 程序时，图形管线功能或多或少是 “关闭” 的

OpenCL 与 CUDA 的关系

OpenCL 是 CUDA 的开放标准版本

CUDA 只能在 NVIDIA GPU 上运行，而 OpenCL 可以在来自许多供应商的 CPU 和 GPU 上运行

在许多方面，OpenCL 与 CUDA 的描述是相似的

因此, Nvidia Tesla 带来的意义是:

虽然应用程序仍然需要通过驱动程序（无论是图形驱动还是计算运行时）才能最终与 GPU 硬件通信，但高级语言或应用程序现在可以通过专用的“计算模式”接口和 API（如 CUDA）来表达和启动程序，绕开了严格的图形渲染管线限制

CUDA Programming Abstractions

CUDA 线程管理

(1) Def: Thread Block | Thread

1. ThreadID, BlockID 的概念, 解释见上图
2. 注意到这个例子的 grid尺寸 是正好可以整除 thread block尺寸 的. 我们将在下面 (4) 介绍“不能整除”应该如何处理

(2) Basic CUDA Syntax

这里以一个经典的 MatrixAdd() CUDA代码进行解析:

• Host Code 顺序执行, 运行在CPU上
• CUDA Device Code 高并发执行, 运行在GPU上

Host Code 和 CUDA Device Code 是泾渭分明的!

(3) CUDA是"显式"线程管理

SPMD “CUDA 线程”的数量在程序中是显式定义的:

线程数量不是自动推断出来的，而是通过主机端 (Host code，即 CPU 运行的代码) 的内核启动语法 (例如 matrixAdd<<<numBlocks, threadsPerBlock>>>)来指定的

(4) 内核调用次数不由数据集合的大小决定

在 CUDA 中, 程序员显式启动一定数量的线程 (例如 72 个线程)

这个数量可能大于或小于程序所需处理的数据集合的大小 (例如, 一个 11×5 的矩阵有 55 个元素, 但程序启动了 72 个线程)

如果线程数量超过数据大小 (例如 11×5 的矩阵启动了 72 个线程), 程序员必须在内核内部通过边界检查 (if (i < Nx && j < Ny) 来保护数组访问不会越界

CUDA Device Memory Model

Host 与 CUDA Device 的 memory 是完全分开的, 不能直接从 Host 访问 CUDA memory

CUDA Device 内部的 Memory 也是逐级的

• Per-thread Private Memory: 每个TB里的每个Thread, 私有
• Per-block Shared Memory: 每个TB里的所有Thread, 共享
• Device Global Memory: 整个Device里的所有Thread, 共享

很明显, 这种"分级"的设计可以体现 "locality"

我们用一个程序简单说一下, global 和 shared 对于 locality 的区别:

CUDA 的其他语法

(1) 类 barrier / sync 的 primitive 作用范围仅限 "同一个TB里的thread"

__syncthreads(): Barrier: wait for all threads in the block to arrive at this point

(2) 类 atomic 的 primitive 作用范围可以做到 shared / global 变量

float atomicAdd(float* addr, float amount): both global memory addresses and per-block shared memory addresses

CUDA Implementation on Modern GPUs

我们经常会在网上看见各种各样的视频, 提到GPU的特点就是core特别多, 但每个core的单核性能远比不上CPU

把CPU比作“独立作战的猛虎”, GPU比作“擅长合作的群狼”

的确, GPU能给支持超大量的“群体行动”, 我们以一个现实的程序来展示一下, 它生成了超过 1 million 个 CUDA threads:

为什么GPU能支持这么多呢？下面我们来讲讲GPU的架构设计!

Compilation and Assignment

(1) 一个编译好的 CUDA Device Binary 包含:

• 程序指令
• metadata
  • block_size
  • local data allocation per thread
  • shared space allocation per TB
  • ...

(2) CUDA 任务调度的最细粒度是 Thread Block:

调度流程:

1. device binary 被划分成众多 "thread block"
   • TB 是一个 "programming-level" 概念, not "hardware-implementation"
   • 可以将 TB 理解成 "work", core 理解成 "workers"
2. 通过 "thread block scheduler" 将TB调度给各个core

重点内容:

• 不同TB之间是完全独立的, CUDA执行它们是并发的, 能以任意顺序执行
• "将TB调度给各个core" 遵循 dynamic shceduling policy, 因此确保维护了资源需求

V100 SM "sub-core"

这一小节, 我们以 Nvidia V100 SM 其中一个 core 内部来进行解析

(1) 区分: 抽象与实现

1. Warp: 是一个 "hardware-implementation", 而不是 "programming-level" 概念
2. Thread Block (TB): 是一个 "programming-level" 概念

(2) 一个sub-core:

• 具有16个warp
• 每个warp: 是32个threads
• 每个thread: 有一个scalar register 和 一个PC
• 每个scalar register中的 \(R_i\) 表示: 单个 CUDA 线程拥有的标量寄存器

(3) 隐式 SIMD:

一个warp包含32个threads, 如果这32个thread都运行的是同一条指令, 那么这个warp就会自动开启SIMD

听起来很智能, 但其实现方式非常简单, 检查这32个thread各自的PC, 看是否相同即可

(4) 指令执行顺序:

这里给个流水线的例子

V100 SM

(1) SM 的组成:

SM: Streaming Multi-processor unit

一个SM包含4个sub-core, 外加 "shared-memory"

(2) 运行方式:

任一运行的”瞬间时刻”, 每个Warp Selector只能选择自己内部的1个运行 (1/16 candidate)

因为一共有4个sub-core, 因此从全局视角来看:

任意时刻, 共有4个warps在运行, 每个sub-core上一个

下面给一个例子, 非常值得研究推敲, 尤其是 "__shared__内存分配" + "warp运行选择"

明明一个SM有64个, 每次只用4个为什么不是一种浪费?

正如上述:

单周期执行的 Warp 数量

• 在每个时钟周期内，每个 Sub-core 会从其管理的 16 个可运行 Warp 中 选择一个 Warp
• 因此，在一个时钟周期内，整个 SM 确实只有 4 个 Warp（每个 Sub-core 1 个）被选中并执行其下一条指令

但这种设计完全不浪费资源, 因其重点是能"隐藏延迟":

(1) 回顾前几讲提到的, "bandwidth is bottleneck":

• 传统的 CPU 设计侧重于减少单个线程的执行时间（例如使用深流水线、乱序执行、复杂分支预测等）
  • 如果一个 CPU 核心只有一个线程在运行，一旦发生内存访问延迟（例如等待从主存 DDR 内存中获取数据），核心的执行单元就必须停下来等待数据，导致资源闲置
• GPU 针对的是 数据并行 负载（例如 Shader Programs 或 CUDA Kernel）
  • 这些程序通常涉及大量的内存访问（尤其是访问速度较慢的 设备全局内存 或 HBM 内存）

(2) Warp 切换的优势:

1. 当一个 Warp (例如 Warp A) 执行一条指令，如果这条指令导致了 长延迟操作 (如访问全局内存), Warp A 就会进入等待状态
2. 在下一个时钟周期，Sub-core 上的 Warp Selector 不会等待 Warp A 完成, 而是会立即选择 另一个可运行的 Warp (例如 Warp B) 来执行其指令
3. 通过在 64 个甚至更多的 Warp 之间快速切换（零开销切换），SM 可以确保其大量的执行单元（例如 V100 SM 包含 \(4 \times 16 = 64\) 个 \(fp32\) MUL-ADD ALUs）始终有工作可做

在一个时钟周期内, 虽然只有 4 个 Warp 在发出指令, 但这 4 个 Warp 几乎可以保证 最大化地利用 SM 核心内部的 64 个 SIMD 执行单元 (Warp)

这种设计使得 GPU 可以在面对高延迟的内存操作时，仍然能保持其计算单元的高利用率，从而实现整体的高吞吐量

V100

Running a CUDA Program on a GPU

Summary

(1) 核心梳理:

1. TB的调度和并发性
   1. TB 调度顺序的独立性
      • 系统假设 TB 之间没有依赖关系, 能以任意顺序调度 TB
      • TB 是 "逻辑并发"(logically concurrent) 的
   2. TB的资源释放
      • 当TB中所有threads完成执行时, 该Block的资源 (共享内存分配 / warp 执行上下文)才会被释放. 以供下一个TB使用
2. TB内部的并发与约束
   1. 同一个TB里的threads是并发执行的
   2. SPMD 程序与协作
      • 一个 CUDA TB 本身就是一个 SPMD 程序
      • TB 里的 threads 是 "workers", 可以通过 __syncthreads() 实现协作
   3. 调度约束
      • 当一个 TB 开始执行时, 其所有线程都必须存在, 并且它们的寄存器状态必须被分配
3. CUDA Implementation on Nvidia GPU
   1. Warp的特性:
      • Warp 概念本身 不存在 于 CUDA 编程模型中, 它是 NVIDIA GPU 上的一个重要的实现细节
   2. SM 调度与通信:
      • 一个 TB "对应的"所有 warp 都被调度到同一个 SM 上
      • 这样做是为了允许通过共享内存变量 (shared memory, on chip) 进行高带宽/低延迟的通信

(2) 回顾一下讲了哪些:

1. GPU发展历程: 传统的图形渲染, 到通用计算
2. GPU如何实现通用计算
   • computer-mode API
   • CUDA <-> GPU
3. GPU架构设计分析

(3) 还有哪些没讲:

1. 我们讲了在2007年后, GPU的computer-mode接口进行通用计算. 但其实graphic-pipeline接口仍作保留
   • graphic-pipeline接口的实现细节/操作方式, 本门课程不做介绍, 见图形学有关课程
2. 我们并没有提到GPU是如何服务于 Deep Learning 的
   • 后面课程会说