GPU架构

一、Roofline 性能分析模型

1.1 问题的提出：计算与访存的瓶颈

当设计一个 CPU 或 GPU 时，衡量其实际性能不能只看浮点计算能力。一个处理器的实际吞吐量受两个因素共同制约：

• 计算能力：每秒能执行多少浮点操作（FLOPS）
• 存储带宽：每秒能从主存中读取多少字节数据

如果数据读不进来，计算单元再快也只能等待。反过来，如果带宽足够但算力跟不上，也无法充分发挥硬件性能。

1.2 运算密度（Arithmetic Intensity）

Roofline 模型（Roofline Model）的核心概念是运算密度（Arithmetic Intensity）：

通俗理解：从内存中拉一个字节数据进来，你对它做了多少次浮点运算。

• 如果拉一个字节做了 3 次浮点运算，运算密度就是 3
• 如果拉两个字节只做了一次运算，运算密度就是 0.5

这个图之所以叫”roofline”（屋檐模型），就是因为它的形状像房屋的屋顶——先是一个斜坡上升，到了顶点以后就平坦了。在斜坡区域，性能受带宽限制（数据读不进来）；在平坦区域，性能受计算能力限制（算力已经拉满，再多的带宽也提不上去）。

1.3 Roofline 曲线的含义

在 Roofline 坐标系中：

• 横轴：运算密度（FLOPs / Byte）
• 纵轴：可达到的浮点性能（GFLOPS / 秒）

曲线的形态：

1. 上升段（带宽受限区）：运算密度较低时，每秒钟能处理多少浮点运算受到数据传入速度的限制。此时提升运算密度能直接带来性能提升。
2. 平坦段（计算受限区）：当运算密度足够大之后，性能达到计算能力的物理上限，不再增长。此时瓶颈在于 ALU 算力。

设计的 CPU 或 GPU 如果落在 roofline 曲线下方，说明要么带宽没达到算法要求，要么计算力没达到算法要求——总而言之是做得不够好。落在曲线上，才说明把硬件的实力发挥到了最大。

1.4 实际应用：神经网络中的运算密度

在深度学习中，运算密度又被称为权重重用（Weight Reuse）：

• 拉入一个权重字节，需要做几次乘法和加法？
• 稠密矩阵（如全连接层）：运算密度较高
• 稀疏矩阵：拉来的数据只做 O(1) 次操作，运算密度低
• FFT 等算法：运算密度为 $O(\log n)$

1.5 能效比（Performance per Watt）

在实际部署中，每瓦电能能产出多少算力至关重要。运行 AI 计算中心受限于电力供应，电费直接决定运营成本。TPU 等专用芯片相对 GPU/CPU 在能效比上有显著优势。

厂商提供的 benchmark 数据需要谨慎解读。公开测试集可能被编译器、运行时或硬件配置针对性优化，真实用户体验往往与跑分存在差距。

二、GPU 体系结构

2.1 CPU 与 GPU 的设计哲学对比

CPU 是把晶体管花在控制逻辑和 cache 上，让一条指令一条指令地跑得飞快。GPU 是把晶体管全部堆在 ALU 上，主打一个”量大”——一幅图像处理完就扔掉，下一帧来了继续，不需要复杂的 control 和巨大的 cache。

2.2 GPU 的层级组织

GPU 内部的层级结构类似于军队编制：

1
2
3
4
5
GPU (整体)
  └── TPC (Thread Processing Cluster) × N
      └── SM (Streaming Multiprocessor) × 3 (per TPC)
          └── SP (Streaming Processor) × 8 (per SM)
              └── ALU (FP Unit + INT Unit)

• TPC：Thread Processing Cluster，线程处理集群
• SM：Streaming Multiprocessor，流式多处理器，是 GPU 的基本调度单元
• SP：Streaming Processor，流式处理器，实际执行计算的单元

以 NVIDIA Tesla C1060（约 15 年前）为例：240 个 core，Core clock: 1.296 GHz，On-board memory: 4GB，Memory IO: 512-bit, 800MHz GDDR3。

GPU 架构就是”套娃”，一环套一环，但没有套出什么高精度的复杂逻辑。这些 SM、SP 之间的通信只能通过 interconnect network（互连网络）来进行，本质上是通过共享 memory 来实现协作。

2.3 SM 内部的 SP 连接

在一个 SM 内部有 8 个 SP。如果所有 SP 都通过共享 memory 来通信，意味着 8 个单元同时访问同一块 memory——总线仲裁的开销会非常大。

直观的想法：让 SP 之间直接通过寄存器互联。

• 全连接：每个 SP 和另外 7 个都连线——图太复杂，画不下
• 分组连接：分成两组，同一分组内的 SP 在一侧，对面的在另一侧，通过交叉连接通信
• 编译器负责把数据交互频繁的线程分配到能直接通信的 SP 上，减少跨越 memory 的访问

在芯片内部这个尺度上，cache 和 memory 本质上是用相同的晶体管设计的，叫什么都一样。关键在于数据通路的组织方式——寄存器间的直连是最快的通信方式。

2.4 GPU 上的并行层次

回顾在 CPU 上学过的并行方式：

1. 流水线并行（Pipeline）：指令级流水
2. 指令级并行（ILP）：多条独立指令同时发射执行
3. 数据级并行（DLP / SIMD）：一条指令同时处理多个数据
4. 线程级并行（TLP / MIMD）：多线程各自执行不同任务

GPU 在这些基础上发展出了 SIMT（Single Instruction Multiple Thread）模型。

三、SIMT 执行模型

3.1 从循环到线程：SIMT 的核心思想

在 GPU 编程中，一个 for 循环的每次迭代天然独立（如 C[i] = A[i] + B[i]），因此可以将每个迭代变成一个线程：

1
2
3
4
5
6
7
8
// CPU 上：一个 for 循环
for (int i = 0; i < N; i++) {
    C[i] = A[i] + B[i];
}

// GPU 上：N 个线程，每个线程执行相同的代码
int tid = blockIdx.x * blockDim.x + threadIdx.x;
C[tid] = A[tid] + B[tid];

把 CPU 线程的传统概念扔掉！在 GPU 里，一个线程的代码可能就是 10-20 行，就是一个循环体而已。一万个线程，每个线程的代码完全一样，只是处理的数据不一样。

3.2 Warp：硬件动态分组的线程束

Warp 的定义：一组执行相同指令（same PC）的线程，由硬件动态划分为一个 Warp。

• 一个 Warp 通常包含 32 个线程
• Warp 内的所有线程在同一时刻执行同一条指令（PC 相同）
• 只要线程们还走在同一行代码上，它们就在同一个 Warp 里

GPU 会根据线程当前执行位置，把位于相同控制流路径上的线程动态组织成 Warp，使它们以相同 PC 执行同一条指令。

3.3 SIMT vs SIMD

SIMT 的全称是 Single Instruction Multiple Thread，也叫 SPMD（Single Program Multiple Data）在 SIMD 硬件上的映射。

SIMT 的关键好处有两个——第一，把代码拆成很多独立的线程，大家一人一个 PC，互不影响；第二，因为灵活了，可以随时把同路的线程重新打包成 Warp 一起干活。就像”先分家过日子，再随时根据需要组团”。

3.4 为什么需要这么多线程

GPU 的线程切换极其轻量——理论上一个 cycle 就能完成切换（CPU 需要约 1000 个 cycle）。但这也意味着：线程数必须大于 1000 时，GPU 才能被有效利用。如果只有几十个线程在 GPU 上跑，频繁的切换开销会严重拖慢性能。

比较：

• CPU：起 4-5 个线程就差不多了，写太多控制不了
• GPU：轻轻松松几千几万个线程，主打一个”量大管饱”

3.5 线程的层级组织（CUDA 中的 thread/block/grid）

CUDA 将线程组织为三级：

1
2
3
Grid（一组 block）
  └── Block（互相协作的线程组）
      └── Thread（基本执行单元）

• Thread：最基本的执行单元，代码只有十几二十行
• Block：可以互相通信、同步的线程组，同一 block 内的线程可以通过 shared memory 共享数据
• Grid：多个 block 组成的整体。不同 block 之间不能直接通信，只能通过 global memory 交换数据

Block 之间有共享 memory，Grid 之间通过 global memory。这就像领导常说的——”有问题先内部沟通、内部解决，实在不行了再上报”。因为上报一次代价非常大（访问 global memory 很慢）。

四、CUDA 软件生态

4.1 NVIDIA 的垄断优势

NVIDIA 真正不可替代的竞争壁垒不是它的 GPU 硬件（硬件谁都可以生产），而是 CUDA 软件生态：

• CUDA 驱动（Driver）：底层，面向硬件
• CUDA 运行时（Runtime）：类似 Java 虚拟机的中间层
• CUDA 库（Libraries）：cuBLAS、cuDNN 等上层数学库

国产 GPU 最大的问题不是芯片做不出来，而是用户用惯了 CUDA，让人家重新用你的编程框架写代码——”我一个普通用户，凭什么让我干这种事？”

4.2 二进制翻译（Binary Translation）

为了使 CUDA 程序能在国产 GPU 上运行，一种直接的方案是：

1. 用户仍然用 CUDA 编写和编译代码
2. 编译出的二进制代码本来只能在 NVIDIA GPU 上执行
3. GPU 厂商在底层做一个二进制翻译器，将 NVIDIA GPU 的二进制指令翻译成自己 GPU 的指令
4. 用户无感知，仍然跑原来的 CUDA 代码

这个想法很直接，在 CPU 领域已经用得很成熟了——国产 CPU 要运行 Windows 的应用程序，很多也走二进制翻译这条路。

但 NVIDIA 从 2024 年开始，在知识产权条款中明确禁止了对 CUDA 程序做二进制翻译——”你用我的环境编译，但不买我的 GPU——没毛病，但不许这么干。”

4.3 CUDA 编程的基本模式

1
2
3
4
5
6
CPU 端（Host）              GPU 端（Device）
────────────                ────────────
1. 分配 GPU 内存
2. CPU→GPU 数据拷贝   →     3. 启动 kernel（成千上万线程并发）
4. GPU→CPU 结果拷贝   ←
5. 释放 GPU 内存

原来的 for 循环：

1
2
for (int i = 0; i < N; i++)
    C[i] = A[i] + B[i];

变成 CUDA kernel 后，把 i 替换为 threadIdx.x + blockIdx.x * blockDim.x，循环被”展开”为 N 个线程并行执行。

五、CPU-GPU 统一内存寻址

5.1 独立 vs 统一地址空间

当多块 GPU 各自拥有独立显存时，每个 GPU 的地址空间都从 0x00 到 0xFF…FF，编程时就需要显式管理：数据在哪个 GPU 上？CPU 如何给指定 GPU 传数据？GPU 之间如何交换数据？

从程序员的角度，统一地址空间最简单——不管数据放在 CPU 还是 GPU，只管申请内存、填指针、做 memory copy，底层怎么走不关心。但分离的地址空间虽然对系统设计者来说最”直接”，却给编程带来了极大的心智负担。

5.2 统一虚拟寻址（Unified Virtual Addressing）

2013 年，NVIDIA 宣布统一内存（Unified Memory）有”光明的未来”（a bright future of unified memory）。

核心思想：CPU 和 GPU 共享同一个 64 位的虚拟地址空间。

• CPU 的 16GB 内存和 GPU 的 16GB 显存在同一个地址空间中各占一段
• 程序员只需调用 cudaMallocManaged() 分配统一内存
• 数据迁移由驱动和硬件自动完成

2013 年的时候大家说”有个 bright future”，十多年后回头看，现在所有买到的 GPU 都是这个 bright future 了。CPU 的虚拟地址空间都搞了多少年了，GPU 这边确实晚了很久，但一旦做出来，就让人觉得”早干什么去了”——做出来之后多有趣。

5.3 GPU 内存层级

等真做设计的时候，这些需求会一个一个自己冒出来——一开始忘了放共享 memory，写着写着发现不对，”来，加一块共享 memory”；后来又发现大家都需要跨组访问，”再加个 global memory”。硬件设计的需求就是这样一层一层自然涌现的。

六、GPU 互连：PCIe 与 NVLink

6.1 PCIe 的局限性

在一台服务器上，CPU 通过 PCIe 总线连接多块 GPU。PCIe 是一个通用总线：可以插 GPU（高速设备），也可以插声卡、打印机控制器（低速设备）。为了兼容各种速度的设备，PCIe 协议设计得比较复杂，这反过来限制了它在高速设备上的性能。

6.2 NVLink：NVIDIA 的高速互连

NVIDIA 设计了自己的专用互连协议 NVLink：不需要兼容低速设备，专门为高速 GPU-GPU 通信设计，速度比 PCIe 快很多。每一个想法都”不过分”——GPU 之间拉一条直连线不过分吧？拉四条也不过份吧？用 NVLink 连接 CPU 和 GPU 也不过份吧？这些需求都是实际使用中自然产生的。

6.3 多 GPU 的连接拓扑

4 个 GPU 的情况：每个 GPU 有 4 个 NVLink 端口。可以做到全连接（每个 GPU 与其他 3 个 GPU 都直连），同时留一个端口连接 CPU。拓扑越密集，通信越快，但成本越高。

8 个 GPU 的情况：每个 GPU 只有 4 个端口，必须设计拓扑，无法做到全互联。矩阵型连接：GPU 排列成网格。重要原则：尽量减少”跳数”（hop count）——从 GPU A 到 GPU B 最好只需一跳，不要经过中间 GPU 转发。

现在大模型时代，”一张嘴就是 1 万个 GPU”。1 万个 GPU 怎么连？这就是专门一节”互连网络”要讨论的内容。

七、SIMT 中的分支处理

7.1 分支发散问题

在 SIMT 模型中，所有线程执行相同的程序。但当程序中有 if-else 分支时，不同线程可能走向不同路径：

1
2
3
4
if (x[tid] != 0)
    y[tid] = x[tid] - y[tid];  // 路径 A
else
    y[tid] = z[tid];           // 路径 B

对于一个 32 线程的 Warp：某些线程走路径 A，另一些走路径 B——PC 不再相同，Warp 分裂了。

7.2 通过 Mask（谓词/屏蔽位）处理分支

GPU 处理分支的方法是：

1. 汇集阶段：分支前，所有线程 PC 相同，是一个完整 Warp
2. 分支阶段：PC 不同了，按照分支方向重新组合 Warp
3. 执行阶段：在每个分支路径内部，线程又形成新的 Warp，用 mask（屏蔽位）标记该路径中哪些线程是活跃的
4. 汇聚阶段：分支结束后，控制流重新汇聚，所有线程再次合并为统一的 Warp

7.3 SIMT Stack

GPU 使用 SIMT 栈（SIMT Stack）来管理分支：

• Control Flow Stack（控制流栈）：保存分支后的返回地址（next PC）
• Active Mask（活跃屏蔽字）：记录当前路径中有哪些线程是活跃的
• 进入分支时 push，退出分支时 pop
• 最终所有路径汇聚到同一条指令时，所有线程重新合并

分支不能预测怎么办？那就别预测了，不要硬来。GPU 的策略是——谁走 A 路就一起走 A，谁走 B 路就一起走 B，到汇聚点大家重新汇合。本质上就是不停地”分家-重组-再分家-再重组”。

八、Fermi 架构要点

Fermi（费米）架构是 NVIDIA 在 GPU 发展史上的重要里程碑，主要创新包括：

• SM 内部包含 32 个 core，分为两组（便于内部管理和共享 memory）
• 快速双精度浮点（Fast Double-Precision FP）
• GPU 存储器 64 位寻址和统一地址空间（与 CPU 统一地址空间）——这是 Fermi 架构最被称道的特性
• L1 cache 64KB 可配置（可部分配置为 shared memory）
• 支持 ECC 内存（错误校验）
• 每 SM 有 4 个 Special Function Unit（特殊函数单元）

• Previous
  环同态
• Next
  事务与并发控制