← 返回笔记列表
🖥️ 硬件基础

GPU 硬件基础：从零理解 CPU / GPU / 寄存器 / SRAM / HBM / PCIe

FlashAttention、TurboQuant 等 LLM 优化论文的硬件基础 · 用图和具体数字把每层讲清楚

🌱
零、从零开始：程序、CPU、线程是什么？

读 FlashAttention、KV Cache 这类文章，你会反复碰到「线程」「核心」「寄存器」「内存」。这一节从头讲清楚这些概念，不需要任何硬件背景。

1. 程序是怎么运行的？

你写了一行 Python：c = a * b，计算机到底做了什么？

图1：一条指令 c = a × b 的执行路径
① 内存 a = 3.14 b = 2.71 取数 ② 寄存器 R1 = 3.14 R2 = 2.71 （紧挨运算器的格子） 运算 ③ 运算单元 R1 × R2 纯电路，1 cycle 写回 ④ 结果 R3 = 8.51 （留在寄存器或写回内存）

所以「内存→寄存器→运算→结果」是所有计算的基本流程。寄存器就是运算单元手边的小格子，数字在里面才能被运算器直接操作。

2. CPU 是什么？

CPU（Central Processing Unit，中央处理器）就是一块能按顺序执行指令的芯片。你可以把它理解为：

🧑‍🍳 CPU = 一个或几个聪明的厨师
每个「核心」（Core）相当于一个厨师，能处理复杂的任务（条件判断、循环、函数调用…），
现代 CPU 有 4–128 个核，每个核一次只做一件事，但做得非常快、非常灵活。

CPU 很快的原因之一是有很大的「L1/L2/L3 Cache」（缓存），这和 GPU 的 SRAM 是类似的概念——把最近用过的数据缓存在芯片上，减少去内存取数的次数。

3. 线程是什么？

「线程」（Thread）是程序执行的最小单位——你可以把它理解为：

👷 线程 = 一个独立的工人
每个线程有自己的「当前在执行哪行代码」（程序计数器）和「私有变量」（寄存器/栈）。
多个线程可以同时运行——比如浏览器打开网页时，一个线程下载 HTML，另一个线程渲染页面，互不干扰。
图2：CPU 核心 vs GPU 线程数量对比（A100）
CPU（如 Intel Xeon） 核心 1 复杂逻辑 核心 2 复杂逻辑 核心 3 复杂逻辑 核心 4 复杂逻辑 GPU（A100：6912 个 CUDA Core） ⋯ 每个小格子 = 1 个简单运算单元 6912 个，同时干同一件事 4 核 · 每核灵活复杂 6912 核 · 每核简单，靠数量取胜

4. GPU 的线程怎么组织的？

GPU 里线程是分级管理的，从小到大依次是：

图3：GPU 线程三层组织结构
Grid（一次 kernel 启动的所有线程） Block（线程块，共享 SRAM） Warp（32 线程） t0 t1 t2 t3 … t28 t29 t30 t31 同时执行同一条指令 Warp（32 线程） t32 t33 … t63 同时执行同一条指令 Block（另一个线程块） Warp（32 线程） 各自私有寄存器 Warp（32 线程） 各自私有寄存器 同一 Block 内的线程共享 SRAM；同一 Warp 内的线程强制同步执行同一条指令
🧵 线程（Thread）
最小执行单元。有自己的寄存器（私有变量）。GPU 上可以有几十万个。
🔗 Warp（线程束）
32 个线程捆绑在一起，强制同步——同一时刻必须执行同一条指令，是 GPU 调度的最小单位。
📦 Block（线程块）
多个 Warp 组成一个 Block。同一 Block 内的线程共享 SRAM，可以通过 SRAM 交换数据。

5. 为什么 GPU 需要这么多线程？

核心原因：隐藏内存延迟
从 HBM 取一个数需要 ~600 cycle（相当于 CPU 等 600 个时钟周期什么都不做）。 GPU 的解法是：准备几千个 Warp，当 Warp A 在等内存时，调度器立刻切换去执行 Warp B、C、D…… 等 A 的数据到了，再切回来继续。靠「轮流等待、从不空闲」来填满这 600 cycle 的空档。
一句话总结： CPU 是赛车手（少量、高速、灵活），GPU 是运动会（大量、同步、简单重复）。 LLM 矩阵运算的特点正好是「大量简单重复」——把同一个乘法做几百亿次——所以 GPU 天然适合。 后面所有的优化（FlashAttention、KV Cache、量化）都是在「怎么更好地喂饱这几千个运算单元」这个问题上做文章。
🔬
零·五、完整图解：一次加法走完所有步骤

上一节讲了概念，这里用一个极简但真实的例子走完全程： C[i] = A[i] + B[i]，计算 4096 个元素的向量加法。 这是 GPU 上最典型的操作，理解了这个，后面 FlashAttention 的数据流就都清楚了。

第一步：你写了一段 Python / CUDA 代码

# Python（PyTorch）——GPU 自动处理并行
C = A + B   # A, B, C 都是 shape=(4096,) 的 tensor，在 GPU 显存里

# 背后 CUDA kernel 做的事（伪代码）：
__global__ void vector_add(float* A, float* B, float* C) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;  # 每个线程知道自己的编号
    C[i] = A[i] + B[i];                            # 每个线程只负责一个位置
}

关键思路：4096 个元素，GPU 同时启动 4096 个线程，每个线程算一个。不是一个接一个算，是全部同时算。

第二步：GPU 把线程按层级组织起来

图1：4096 个线程如何分组（以每 Block 128 线程为例）
Grid = 整个任务（4096 个线程） Block 0（线程 0~127，共 128 个，共享 SRAM） Warp 0（线程 0~31） t0 t1 t2 t3 … t8 t9 t10 t11 … （共 32 个线程） 同一时刻执行 同一条指令！ Warp 1（线程 32~63） （共 32 个线程） 也在同时 执行同一条指令 Block 1（线程 128~255，共享另一块 SRAM） Warp 4（线程 128~159） 32 个线程 Warp 5（线程 160~191） 32 个线程 … （共 32 个 Block，每块 4 个 Warp）

第三步：SM 是什么？（它才是 GPU 里真正的「核」）

图2：SM 内部结构（以一个 SM 为例）
SM（流式多处理器）= GPU 里的一个「核」 A100 有 108 个 SM，每个 SM 包含以下组件 寄存器文件（Register File） 整个 SM 共享一个大池子（A100：256KB/SM） 线程启动时，每个线程从池子里分到一份（如 32~64 个寄存器） ← 不是「每个运算单元固定配几个」，而是动态分配 CUDA Core（运算单元）× 64 个（A100 每 SM） + + + + + + + + … （共 64 个） 每个 CUDA Core 是最基本的加减乘除单元，同一时刻只能做一件事（加法 OR 乘法） SRAM（共享缓存，192KB/SM） 同一 Block 内所有线程都可访问 用于线程间共享中间结果（如 FlashAttention 的 Q/K tile） ~20-30 cycle 延迟，远快于 HBM Warp 调度器（Warp Scheduler） 监控哪些 Warp 的数据已就绪 → 优先执行就绪的 Warp → 当前 Warp 等内存时立即切换到下一个就绪 Warp
📐 层级关系总结（回答你的问题）
线程（Thread）
最小执行单元，有自己的寄存器（从 SM 寄存器文件里分到一份）
Warp（32个线程）
GPU 调度的最小单位，同一 Warp 内 32 个线程强制同步，执行同一条指令
Block（多个Warp）
同一 Block 内的线程共享 SRAM，可以通过 SRAM 交换数据；Block 分配到某个 SM 上运行
SM（多个Block）
GPU 里真正的「核」，相当于一个强大的 CPU 核心，有自己的寄存器文件、SRAM、运算单元、调度器。A100 有 108 个 SM。
GPU（多个SM）
整个 GPU = 108 个 SM + 共享的 HBM（显存），通过 PCIe 连接到 CPU

第四步：线程 #5 计算 C[5] = A[5] + B[5] 的完整旅程

以线程 #5（在 Warp 0，Block 0，被调度到 SM #3）为例，走完所有步骤：

图3：线程 #5 计算 C[5] = A[5] + B[5] 的完整数据流
← 空间视角（从左到右：数据从远处搬到运算单元） HBM（显存） 80GB · 2TB/s A[5] = 1.2 B[5] = 3.4 C[5] = ? → 4.6 ⚠ ~600 cycle 读写延迟大 ① 读取 ~600 cy SM #3（被分配到这个 SM 上运行） SRAM（片上缓存） A[5]=1.2 B[5]=3.4 （对于简单加法可能直接走寄存器， 但 tile 运算时会先在 SRAM 缓存一批） ② 装入 ~20 cy 线程 #5 的寄存器 R1=1.2 R2=3.4 R3=? 每个线程私有，只有本线程可见 ③ 加法 1 cycle! CUDA Core（加法器） R1 + R2 → R3 1.2 + 3.4 = 4.6（纯电路，1 个 cycle，不碰任何内存！） ④ 写回 C[5]=4.6 → HBM（~600 cycle） 同时！Warp 0 里另外 31 个线程也在各自计算 t0~t4 算 C[0]~C[4]，t6~t31 算 C[6]~C[31] 步骤完全一样，32 个线程强制同步进行

第五步：时间线视角——Warp 调度如何隐藏延迟

刚才我们说第①步需要 600 cycle，第③步只需要 1 cycle。那 600 cycle 里 SM 在干什么？答案是：切换去执行其他 Warp。

图4：SM 时间线——多个 Warp 轮流执行，填满等待空档
时间 → Warp 0 读内存 等待 HBM 数据（~600 cycle，Warp 挂起） 算 写回 Warp 1 等待 读内存 等待 HBM 数据 算 Warp 2 等待 读内存 等待 HBM 数据 SM 实际 在执行: Warp0 Warp1 Warp2 Warp3 Warp4 Warp5 Warp0✓ Warp… → SM 从不空闲！ 关键：SM 只有 1 个调度器，但同时持有几十个 Warp，哪个 Warp 数据就绪就执行哪个，永不等待
💡 两个基础概念：cycle 是什么？加法是怎么「算出来」的？
① cycle（时钟周期）= GPU 工作的最小节拍

CPU/GPU 内部有一个「时钟」，每隔固定时间发出一次「滴答」信号，驱动所有电路同步工作。这个「滴答一次」就是 1 cycle。
A100 GPU 时钟频率 ≈ 1.4 GHz，即每秒 14 亿次滴答，1 cycle ≈ 0.7 纳秒。

不同操作消耗的 cycle 数对比
cycle数 1 加法运算 （纯电路） ~20 从 SRAM 取数 （片上缓存） ~600 从 HBM 取数 （片外显存）
② 加法是怎么「算出来」的？—— 1-bit 加法器原理图

计算机里所有数字都是二进制（只有 0 和 1）。加法电路的核心是一个叫 全加器（Full Adder）的电路，由几个逻辑门（AND / XOR / OR）搭成，不需要任何程序指令，纯粹是电压传播。下面是 1+1=10（二进制）的完整过程：

全加器（Full Adder）分步图解：计算 1 + 1 = ?（二进制）
XOR 门（异或）
两个输入不同时输出 1
两个输入相同时输出 0
0 XOR 0 = 0  |  1 XOR 0 = 1
0 XOR 1 = 1  |  1 XOR 1 = 0
AND 门（与）
两个输入都是 1 才输出 1
否则输出 0
0 AND 0 = 0  |  1 AND 0 = 0
0 AND 1 = 0  |  1 AND 1 = 1
OR 门（或）
至少一个输入是 1 就输出 1
两个都是 0 才输出 0
0 OR 0 = 0  |  1 OR 0 = 1
0 OR 1 = 1  |  1 OR 1 = 1
输入：A=1，B=1，Cin（上一位的进位）=0 步骤①  XOR 门 #1：A XOR B = 本位中间值 A=1 A=1 A = 1 B = 1 XOR 1 XOR 1 = 0 中间值 S1 = 0 （还不是最终答案） 步骤②  AND 门 #1：A AND B = A+B 是否有进位？ A = 1 B = 1 AND 1 AND 1 = 1 进位标志 C1 = 1（有进位！） 步骤③  XOR 门 #2：S1 XOR Cin = 最终本位（Sum） S1 = 0 Cin = 0 XOR 0 XOR 0 = 0 Sum = 0 ← 最终本位结果 步骤④  AND 门 #2 + OR 门：计算最终进位（Cout） S1=0 Cin=0 AND 0 AND 0 = 0 C1=1 OR 0 OR 1 = 1 Cout = 1 ← 最终进位
最终结果： Sum（本位）= 0，Cout（进位）= 1
二进制写法：10 = 十进制 2  ✓ （1 + 1 = 2，正确！）
整个过程就是电压在 5 个逻辑门之间依次传播，没有任何程序指令，1 个 cycle 内完成。

为什么叫「纯电路」？因为加法的答案不是「查表」「执行指令」得到的，而是电压在导线和晶体管里传播，到达输出端时物理上就是正确答案。整个 A100 GPU 同时有数千个这样的加法器在工作。

32-bit 浮点数的加法器比 1-bit 版本复杂得多（需要处理符号位、指数位、尾数位），但原理一样：都是逻辑门的组合电路，不经过任何「软件」。

③ 寄存器是什么？——从晶体管到寄存器的层级

你问的「寄存器是不是多个二极管」——方向对，但材料是晶体管（不是二极管）。层级关系如下：

晶体管 可当电子开关 通/断 = 1/0 组合 逻辑门 AND / OR / XOR / NOT 每个门 ~6 个晶体管 交叉连接 触发器（Flip-Flop） 能记住 1 个 bit（0 或 1） ~6 个逻辑门，~36 个晶体管 N 个并排 寄存器 N 个触发器并排 存 N-bit 数字 例：十进制 4 = 二进制 100，需要 3 个触发器并排（3-bit 寄存器） 1 0 0 ← 3 个触发器，各存 1 bit，合起来就是 4 bit 2 触发器① 触发器② 触发器③ GPU 的 32-bit 寄存器 = 32 个触发器并排 32 个触发器 × ~36 个晶体管/个 ≈ 1152 个晶体管，只存一个浮点数（如 1.2） A100 一个 SM 有 256KB 寄存器文件 ≈ 200 亿个晶体管，全在 GPU 芯片上

汇总：各步骤时间成本对比

步骤做什么延迟谁在等
① 从 HBM 读数据 A[5], B[5] → 进入 SM（寄存器或 SRAM） ~600 cycle Warp 0 挂起，SM 去执行其他 Warp
② SRAM → 寄存器 从片上缓存装入线程私有寄存器 ~20 cycle 极短，几乎可忽略
③ 加法运算 R1 + R2 → R3，纯电路（见下方解释） 1 cycle ⚡ 无等待，不碰任何存储
④ 写回 HBM R3（结果）→ HBM C[5] ~600 cycle 同样 Warp 挂起，SM 去执行其他 Warp
整个流程的核心洞察：
真正的计算（第③步）只要 1 cycle，但数据搬运（第①④步）各需要 600 cycle。 优化 GPU 性能 = 减少①④的次数，或者用多 Warp 轮流来填满①④的等待时间。 这就是为什么 FlashAttention 的核心改进是「减少 HBM 读写次数」而不是「让乘法更快」。
🗺️
一、整体架构全景

理解 LLM 推理优化（FlashAttention、KV Cache 量化等）首先需要搞清楚一件事：数据存在哪里，如何到达计算单元。

主机（CPU 侧） CPU 控制逻辑 · PyTorch 主程序 L1/L2/L3 Cache MB 级 · CPU 片上 主机内存（DRAM） 16–512 GB · 50–100 GB/s 模型文件先存这里 训练前 .to('cuda') 复制到 GPU 操作系统、Python、 用户数据均在此 PCIe ~32 GB/s 最慢环节！ GPU（如 NVIDIA A100 / H100） HBM（显存） High Bandwidth Memory 40–80 GB 带宽 ~2 TB/s 模型权重 KV Cache 激活值 / 中间结果 梯度（训练时） ⚠ 离计算单元有距离 读写有延迟，是主要瓶颈 读 → ← 写 SM（流式多处理器） A100: 108 个 SM · H100: 132 个 SRAM（片上缓存） L1 Cache + Shared Memory 每 SM ~256 KB · 全 GPU ~20 MB 带宽 ~10–20 TB/s ⚡ 容量小，但极快，是 FlashAttn 的关键 寄存器 每线程私有 ~256 个 32-bit 最快！直接接运算器 Tensor Core（矩阵乘） CUDA Core（通用运算） FlashAttention：让 S/P 只在 SRAM 和寄存器中生存，永不写 HBM → HBM IO 减少 5–10×
一句话总览：GPU 是一台专门做矩阵运算的小型计算机，有自己的「CPU」（叫 SM）、自己的「L1 Cache」（叫 SRAM）、自己的「内存」（叫 HBM），通过 PCIe 连接到主机 CPU。LLM 推理的所有优化，核心都是「减少慢速 HBM 的读写次数，让运算尽量在快速 SRAM 和寄存器里完成」。
📦
二、五层存储层级（从最快到最慢）
层级 名称 容量 带宽 / 延迟 位置 类比
① 寄存器 Register File 每线程 ~256 个 32-bit 无限快（直接接运算器） SM 内部，每线程私有 手里正拿着的东西
② SRAM 片上 L1 Cache / Shared Memory 每 SM ~256 KB；全 GPU ~20 MB 10–20 TB/s · <5 ns SM 内部，同 SM 线程共享 工作台
③ HBM High Bandwidth Memory（显存） 40–80 GB ~2 TB/s · ~100 ns GPU 封装内，所有 SM 共享 GPU 的仓库
④ 主机内存 DRAM / RAM 16–512 GB 50–100 GB/s · ~100 ns CPU 主板上 CPU 的仓库
⑤ 磁盘 SSD / NVMe TB 级 2–10 GB/s · ~100 μs 独立存储设备 档案室
带宽可视化（对数比例）
寄存器
近似无限（电路直连）
SRAM
10–20 TB/s
HBM
~2 TB/s
主机内存
50–100 GB/s
PCIe
~32 GB/s
💡 为什么 HBM 明明有 2 TB/s 带宽，还会是瓶颈？

因为 Tensor Core 算力达到 312 TFLOPS（A100，fp16），即每秒 312 万亿次运算。假设每次运算消耗 2 bytes 数据（fp16），理论上需要 312×2 = 624 TB/s 的带宽才能喂饱算力——但 HBM 只有 2 TB/s，相差 312 倍。

所以 GPU 大量时间在空转等待读数据，而非真正在算。这就是「内存墙（Memory Wall）」。FlashAttention 的本质就是减少无谓的 HBM 读写，让 Tensor Core 少等待。

⚙️
三、GPU 内部结构：SM / CUDA Core / Tensor Core

一块 GPU 不是一个单一的大型计算器，而是由数百个独立的「小型计算单元」组成，每个叫 SM（Streaming Multiprocessor，流式多处理器）。

SM（流式多处理器）
GPU 的基本计算单元，类比 CPU 的一个核。A100 有 108 个 SM，H100 有 132 个。每个 SM 包含：
  • 若干 CUDA Core（通用运算）
  • 若干 Tensor Core（矩阵加速）
  • 专属 SRAM（L1 Cache + Shared Memory）
  • 寄存器文件（每个线程私有）
CUDA Core
通用浮点/整数运算单元，可以做：
  • 加减乘除（标量运算）
  • softmax、激活函数（逐元素）
  • 条件判断、数据类型转换
A100 每个 SM 有 64 个 FP32 CUDA Core（共 6912 个）
Tensor Core
专门为矩阵乘法设计的硬件单元，一次指令可完成 4×4×4 的矩阵乘加（MMA），比 CUDA Core 快约 8 倍。
  • 矩阵乘法 Q@K.T → Tensor Core
  • 线性层 X@W → Tensor Core
  • softmax、ReLU → CUDA Core

GPU vs CPU 的根本差异

CPU — 少量强大核心
A modern server CPU 有 32–128 个大型核心，每个核有复杂的分支预测、乱序执行、深度流水线。

适合：复杂的控制流（if/else/while），单线程延迟敏感任务，不规则数据访问。
GPU — 大量简单核心
GPU 有 6000–18000 个简单 CUDA Core，每个核极其简单，但靠数量取胜。

适合：所有线程做完全相同的操作（如矩阵乘法），数据规则、分支极少的大规模并行计算。
🧵
四、Warp：GPU 的最小调度单位

GPU 里的线程不是单独调度的，而是以 32 个线程为一组，称为一个 Warp。同一个 Warp 里的所有线程在任意时刻都执行 同一条指令（SIMT：Single Instruction Multiple Threads）。

✅ 理想情况：所有线程执行相同操作
如矩阵乘法，每个线程计算输出矩阵的一个元素，操作完全相同。32 个线程同时发出同一条乘加指令，Tensor Core 以最高吞吐量运行。
线程 0: C[0] += A[0] * B[0]
线程 1: C[1] += A[1] * B[1]
……（32线程同时执行同一指令）
❌ Warp Divergence：线程走不同分支
如果 32 个线程里有些走 if、有些走 else，GPU 只能串行处理：
if (x > 0): ← 部分线程走这里
  result = sqrt(x)
else: ← 其余线程等待
  result = 0
↓ 先跑 if 分支，else 线程等待
↓ 再跑 else 分支，if 线程等待
→ 总时间 = if时间 + else时间
这叫 Warp Divergence，有效吞吐量骤降。
这就是 RabitQ 在 GPU 上慢 174 万倍的根本原因
RabitQ 的格点搜索需要 binary search + recursive subdivision，充满条件分支。大量 Warp Divergence 让 GPU 变成「串行执行器」。而 FlashAttention / TurboQuant 的量化/矩阵乘法几乎没有分支，所有线程做完全相同的操作，GPU 满负荷并行。
🧱
五、内存墙（Memory Wall）：为什么带宽是瓶颈

计算机体系结构中，「内存墙」是指 CPU/GPU 的算力增长速度远快于内存带宽的增长速度，导致处理器大量时间在等待数据，而非真正在计算。

具体数字（A100）

指标数值含义
Tensor Core 算力（fp16）312 TFLOPS每秒 3120 亿次 fp16 乘加
HBM 带宽2 TB/s每秒最多读 2 TB 数据
喂饱算力需要的带宽~624 TB/s312T ops × 2B = 624 TB
差距~312 倍算力是带宽的 312 倍
💡 矩阵乘法的算术强度（Arithmetic Intensity）

「算术强度」= 计算量（FLOPs） / 数据量（Bytes），单位是 FLOP/Byte。

A100 的「峰值算术强度」（roofline 拐点）约为：312 TFLOPS / 2 TB/s = 156 FLOP/Byte。

  • 大矩阵乘法（如 m=n=k=4096）：算术强度 ≈ 2×4096 = 8192 FLOP/Byte → 远超 156，Tensor Core 是瓶颈，内存足够
  • 向量加法（C = A + B）：算术强度 = 1 FLOP / 6 Byte ≈ 0.17 → 远低于 156，带宽是瓶颈
  • 标准 Attention（n=4096）：S 矩阵 32 MB，O 矩阵 1 MB，反复读写 HBM → 带宽密集型，是瓶颈

FlashAttention 通过 Tiling 把 Attention 变成「计算密集型」操作，让 Tensor Core 不再等带宽。

🔢
六、寄存器：GPU 里最快的存储是什么？

寄存器是什么？

寄存器（Register）是 GPU 里离运算单元最近、速度最快的存储——事实上它根本不是「内存」， 而是运算电路本身的一部分，是直接连在加法器/乘法器旁边的几个「小格子」， 用来临时存放当前正在参与运算的数字。

图1：一个 CUDA 线程视角下的存储层级
← 越往左越快、越贵、越小 🔢 寄存器 每线程私有 ~255 个/线程 总量 ~16 MB/SM 延迟：~1 cycle 在运算电路内部 溢出 🟢 SRAM L1 Cache + Shared Mem 192 KB/SM（A100） 所有线程共享 延迟：~20-30 cycle 在 SM 芯片上 溢出 🔵 HBM（显存） GPU 全局内存 80 GB（A100） 所有 SM 共享 延迟：~600-700 cycle 在 GPU 封装上（HBM 芯片） PCIe 🟣 CPU 内存 DRAM 延迟：~μs 级 不在 GPU 上

寄存器和 SRAM 有什么区别？

🔢 寄存器
  • 每个线程私有，别的线程看不到
  • 存放「这个线程正在算的中间值」
  • 就在乘法器旁边，0 延迟传输
  • A100 每个 SM 约 256KB 寄存器文件
  • 每线程最多 255 个（超出则 spill 到 SRAM）
🟢 SRAM（Shared Memory）
  • 同一 SM 内所有线程共享
  • 存放「这批线程都要用的数据块」（如 Q/K tile）
  • 需要显式加载（__shared__ 变量）
  • A100 每个 SM 最大 192 KB
  • 比寄存器慢一点，但远快于 HBM

图2：一次矩阵乘法里，寄存器在哪里工作？

图2：计算 c = a × b 的完整数据流
HBM a = 3.14 b = 2.71 （全部权重/激活） ① 读入 ~600 cycle SRAM 缓存 a, b 的 tile （FlashAttention 里 Q tile / K tile 就在这里） A100: 192 KB/SM ② 装入 ~20 cycle SM 内部：寄存器 + 运算器 R1 a = 3.14 R2 b = 2.71 × 纯电路，0延迟 R3 c = 8.51 ③ 运算在寄存器间完成，不触碰任何存储！ ④ 如果结果要持久化，才写回 SRAM 或 HBM
关键理解：第③步（乘法本身）完全不碰内存
乘法电路只操作寄存器里的值——寄存器就在乘法器旁边，相当于「随手一摸就到」。 真正慢的是第①步（从 HBM 把数据搬来），这需要 600+ cycle； 第③步只需要 1 个 cycle，快 600 倍。
所以「优化 GPU 计算」的本质就是：减少去 HBM 取数据的次数。

寄存器溢出（Register Spilling）是什么？

每个线程最多用 255 个寄存器。如果一个线程需要的中间变量太多，放不下，就会「溢出」到 SRAM（更慢），SRAM 也放不下则溢出到 HBM（更更慢）。这叫 register spilling，会显著拖慢 kernel 速度。

图3：寄存器溢出的代价
正常情况
🔢→✖️→🔢
寄存器间运算
~1 cycle
›
溢出到 SRAM
🔢→🟢→🔢
多一次 SRAM 读写
~20 cycle
›
溢出到 HBM
🔢→🔵→🔢
多一次 HBM 读写
~600 cycle

跟 FlashAttention 的关系

✅ FlashAttention：结果小 → 留在寄存器
每次只算一个 tile（如 64×64），tile 内的 S_tile、P_tile 结果只有几 KB，寄存器 + SRAM 放得下， 计算完直接继续，永不写 HBM。
❌ 标准 Attention：结果大 → 必须写 HBM
S 矩阵是 4096×4096 fp16 = 32 MB，寄存器和 SRAM 根本装不下， 必须写到 HBM，下次 softmax 时再从 HBM 读回——往返 2 次，这才是瓶颈。
一句话总结：寄存器是「运算员手里的草稿纸」
HBM 是书架（大但远），SRAM 是桌上的便签本（小但近），寄存器是手里拿着的那张纸（最小但随手就用）。 GPU 优化的终极目标：让运算员尽量只用手里的草稿纸，少去书架取书。
🔌
七、PCIe 与多 GPU 通信

CPU（主机）和 GPU（设备）通过 PCIe（Peripheral Component Interconnect Express）总线 连接。

通信链路带宽典型场景
PCIe 4.0 × 16（CPU ↔ GPU）~32 GB/s 双向模型权重上传、推理结果下载
PCIe 5.0 × 16（新平台）~64 GB/s 双向下一代服务器
NVLink（GPU ↔ GPU，A100）~600 GB/s 双向多 GPU All-Reduce 梯度同步
NVLink（H100）~900 GB/s 双向超大模型张量并行
PyTorch 中 .to('cuda') / .cuda() 就是触发 PCIe 传输
把一个 7B 参数模型（fp16，~14 GB）从 CPU 内存上传到 GPU 显存，通过 PCIe 4.0 至少需要 14 GB / 32 GB/s = ~0.44 秒。这就是为什么模型一般只在启动时加载一次，推理期间权重常驻 GPU。

多 GPU 训练为什么需要 NVLink？

以 8 块 A100 训练一个大模型，每次反向传播需要同步所有 GPU 的梯度（All-Reduce）。假设梯度数据量 14 GB：

  • 走 PCIe：14 GB / 32 GB/s × 8 = 3.5 秒（严重拖慢训练）
  • 走 NVLink：14 GB / 600 GB/s × 8 = 0.19 秒（可接受）

NVLink 让多 GPU 之间的通信带宽接近单 GPU 内部的 HBM 带宽，这是大模型训练能 Scale 到千卡的关键基础设施。

📊
八、关键数字速查表（A100 SXM4 80GB）
指标数值备注
Tensor Core 算力（fp16）312 TFLOPS矩阵乘法峰值
HBM 容量80 GB存模型权重、KV Cache 等
HBM 带宽2 TB/s实际利用率约 70–80%
SRAM 总量~40 MB108 SM × ~256 KB（含 L2 Cache ~40 MB）
SRAM 带宽~10–20 TB/sHBM 的 5–10 倍
SM 数量108每 SM 64 FP32 CUDA Core + 4 Tensor Core
线程总数（最大）221,184108 SM × 2048 线程/SM
Warp 大小32 线程固定，所有 NVIDIA GPU 相同
PCIe 带宽~64 GB/sPCIe 4.0 × 16，双向
NVLink 带宽600 GB/s双向，多 GPU 互联
推荐阅读：这些知识在笔记中的应用
  • LLM 推理优化笔记 — FlashAttention 如何利用 SRAM Tiling，TurboQuant 量化如何降低 HBM 带宽压力