学习笔记

GPU 硬件架构知识

24 分钟阅读

学习定位:从零开始理解 GPU 为什么长这样、每个部件是干嘛的、瓶颈在哪里、写代码时如何利用这些硬件特性。读完本篇,你将建立起"硬件结构 → 性能瓶颈 → 代码优化"的完整认知链路。 预计阅读时间:40 分钟。

一、开篇:为什么需要 GPU?从 CPU 的"困境"说起

1.1 CPU 的设计哲学:追求"快",而非"多"

现代 CPU(如 Intel Core、AMD Ryzen、Apple M 系列)的设计理念是让单个任务执行得尽可能快。它就像一个全能型精英专家

  • 复杂控制逻辑:强大的分支预测、乱序执行、 speculative execution
  • 庞大缓存 hierarchy:L1/L2/L3 Cache 占比很高,为了减少访存延迟
  • 少量核心:消费级 CPU 通常 8~32 核,服务器级也不过 64~128 核
  • 高主频:3~6 GHz,追求每时钟周期做更多事

CPU 的困境:面对 AI 训练这种"海量相同计算"的场景,CPU 就像让 16 个顶级教授去手抄一百万份试卷——教授再厉害,人数太少,抄不完。

1.2 GPU 的设计哲学:追求"多",用数量碾压

GPU 的设计理念截然相反:我不追求单个线程多快,我要同时跑几万个线程

  • 极简控制逻辑:没有复杂分支预测,每个核心非常简单
  • 小缓存:缓存占比低,把晶体管省下来做计算单元
  • 海量核心:H100 有 132 个 SM,每个 SM 里有上百个执行单元,总计相当于上万个"迷你工人"
  • 适中主频:1~2 GHz,但并行度极高

类比

场景CPUGPU
比喻16 个顶级教授10,000 个普通高中生
擅长解一道极难的数学题每人解一道简单的数学题
AI 训练教授们讨论半天10,000 人同时算矩阵乘法

结论:CPU 是**低延迟(Latency-oriented)架构,GPU 是高吞吐(Throughput-oriented)**架构。AI 计算(矩阵乘法)属于"计算高度规则、数据量大、可并行度极高"的任务,天然适合 GPU。


二、GPU 物理架构全景图

2.1 从芯片到板卡

一块你插在服务器里的 NVIDIA GPU 加速卡,内部结构从大到小是这样的:

┌─────────────────────────────────────────────────────────────┐
│                    GPU 加速卡(如 A100/H100)                  │
│  ┌───────────────────────────────────────────────────────┐  │
│  │              GPU 芯片(Die / Silicon)                 │  │
│  │  ┌─────────────────────────────────────────────────┐  │  │
│  │  │           GPC(图形处理集群)× N 个               │  │  │
│  │  │  ┌─────────────────────────────────────────┐    │  │  │
│  │  │  │      TPC(纹理处理集群)× M 个           │    │  │  │
│  │  │  │  ┌─────────────────────────────────┐   │    │  │  │
│  │  │  │  │    SM(流多处理器)× 2 个        │   │    │  │  │
│  │  │  │  │  ┌─────────────────────────┐   │   │    │  │  │
│  │  │  │  │  │  CUDA Core / Tensor Core │   │   │    │  │  │
│  │  │  │  │  │  寄存器 / Shared Memory  │   │   │    │  │  │
│  │  │  │  │  │  Warp Scheduler           │   │   │    │  │  │
│  │  │  │  │  └─────────────────────────┘   │   │    │  │  │
│  │  │  │  └─────────────────────────────────┘   │    │  │  │
│  │  │  └─────────────────────────────────────────┘    │  │  │
│  │  │                                                 │  │  │
│  │  │  L2 Cache(所有 GPC 共享)                       │  │  │
│  │  │  NVLink / PCIe 控制器                           │  │  │
│  │  │  HBM 控制器(连接显存芯片)                       │  │  │
│  │  └─────────────────────────────────────────────────┘  │  │
│  │                                                       │  │
│  │  HBM 显存芯片(2.5D/3D 封装在 GPU 周围)              │  │
│  └───────────────────────────────────────────────────────┘  │
└─────────────────────────────────────────────────────────────┘

2.2 计算单元的层级组织

层级英文全称中文作用数量参考(H100 SXM)
GPCGraphics Processing Cluster图形处理集群一组 TPC 的集合,共享光栅化单元等~8 个
TPCTexture Processing Cluster纹理处理集群包含 2 个 SM,共享纹理采样单元~22 个
SMStreaming Multiprocessor流多处理器GPU 最基本的计算单元,包含 Core、寄存器、Shared Memory132 个

注意:虽然名字里带"图形"(Graphics),但现代数据中心 GPU(A100/H100)里的 GPC/TPC 主要是为了兼容历史命名,实际工作中几乎 100% 用于通用计算(GPGPU)


三、SM(流多处理器):GPU 的计算心脏

SM 是 NVIDIA GPU 最核心的计算单元。理解 SM,就理解了 GPU 80% 的工作原理。

3.1 SM 是什么?一个"微型工厂"

每个 SM 就像一个独立的微型计算工厂,有自己的工人(CUDA Core)、专业设备(Tensor Core)、原材料暂存区(寄存器、Shared Memory)和调度主任(Warp Scheduler)。

H100 的 SM 内部结构(简化版):

┌──────────────────────────────────────────────────────────────┐
│                    Streaming Multiprocessor (SM)              │
│                                                               │
│  ┌──────────────┐ ┌──────────────┐ ┌──────────────────────┐  │
│  │ Warp Scheduler│ │ Warp Scheduler│ │ Warp Scheduler        │  │
│  │ Warp Scheduler│ │ Warp Scheduler│ │  (共 4 个调度器)       │  │
│  └──────┬───────┘ └──────┬───────┘ └──────────┬───────────┘  │
│         │                │                    │              │
│  ┌──────▼───────┐ ┌──────▼───────┐ ┌──────────▼───────────┐  │
│  │ CUDA Core    │ │ CUDA Core    │ │ Tensor Core           │  │
│  │ FP32 × 64    │ │ INT32 × 64   │ │ 4 个(第四代)        │  │
│  │ FP64 × 32    │ │              │ │                       │  │
│  └──────────────┘ └──────────────┘ └──────────────────────┘  │
│                                                               │
│  ┌─────────────────────────────────────────────────────────┐ │
│  │ 寄存器文件(Register File)                              │ │
│  │ 容量:256 KB / SM                                        │ │
│  │ 作用:每个线程的私有高速存储                              │ │
│  └─────────────────────────────────────────────────────────┘ │
│                                                               │
│  ┌─────────────────────────────────────────────────────────┐ │
│  │ L1 Cache / Shared Memory(共享内存)                     │ │
│  │ 容量:256 KB / SM(可编程分配比例)                      │ │
│  │ 作用:同 Block 内线程共享的高速缓存                        │ │
│  └─────────────────────────────────────────────────────────┘ │
│                                                               │
│  ┌─────────────────────────────────────────────────────────┐ │
│  │ Load/Store 单元                                          │ │
│  │ 作用:负责与 L2 Cache / HBM 之间的数据搬运                │ │
│  └─────────────────────────────────────────────────────────┘ │
└──────────────────────────────────────────────────────────────┘

3.2 核心组件详解

3.2.1 CUDA Core:通用计算的"基础工人"

CUDA Core 是 SM 里最基础的计算单元,负责执行标量运算(一次算一个数)和简单的向量运算。

类型作用比喻
FP32 Core单精度浮点运算普通计算器
FP64 Core双精度浮点运算科学计算器(数量通常只有 FP32 的一半)
INT32 Core整数运算整数计算器

关键事实

  • A100 每个 SM 有 64 个 FP32 CUDA Core,108 个 SM,总共约 6,912 个 FP32 Core
  • H100 每个 SM 有 128 个 FP32 CUDA Core,132 个 SM,总共约 16,896 个 FP32 Core
  • CUDA Core 是"全能选手",但处理矩阵乘法的效率远不如 Tensor Core

3.2.2 Tensor Core:矩阵计算的"专业大厨"

Tensor Core 是 NVIDIA 从 Volta 架构(V100,2017 年)开始引入的专用硬件单元,专门加速矩阵乘加运算(Matrix Multiply-Accumulate, MMA)

为什么需要 Tensor Core?

深度学习 90% 的计算都是 C = A × B + C(GEMM)。如果用 CUDA Core 做这个运算,需要三层嵌套循环,每个乘法都要占用一个 CUDA Core 很多个时钟周期。

Tensor Core 的思路是:把矩阵乘法做成一个硬件电路,一拍搞定一个小矩阵块

各代 Tensor Core 演进:

架构代表 GPUTensor Core 代数核心能力
VoltaV100第一代支持 FP16 矩阵运算
TuringRTX 2080 Ti第二代增加 INT8/INT4 支持
AmpereA100第三代增加 TF32、BF16、FP64、稀疏性加速
HopperH100/H200第四代增加 FP8、Transformer Engine、动态精度切换
BlackwellB200第五代FP4/FP6、第二代 Transformer Engine

Tensor Core 支持的精度(以 H100 为例):

数据类型说明典型用途
FP64双精度浮点科学计算(HPC)
FP32单精度浮点通用计算
TF32TensorFloat-32(19 位)A100 引入,训练默认精度,速度接近 FP16,精度接近 FP32
FP16 / BF16半精度训练主力,速度翻倍
FP88 位浮点H100 引入,推理和部分训练,速度再翻倍
INT8 / INT4整数量化推理

初学者误区:Tensor Core 不是取代 CUDA Core,而是各司其职。Tensor Core 只管矩阵乘法,其他事情(如数据预处理、索引计算、条件判断)还是要 CUDA Core 来做。

3.2.3 Warp 与 Warp Scheduler:线程的"班长"和"调度主任"

Warp(线程束)

  • GPU 不以"单个线程"为单位调度,而是以 Warp = 32 个线程 为一组进行调度
  • 一个 Warp 里的 32 个线程执行同一条指令(SIMT:Single Instruction, Multiple Threads)
  • 如果 Warp 内线程发生分支(if/else),会导致线程发散(Divergence),性能下降

Warp Scheduler

  • 每个 SM 有 4 个 Warp Scheduler
  • 每个时钟周期,每个 Scheduler 可以挑选一个就绪的 Warp 发射指令
  • 当某个 Warp 在等待数据(访存延迟),Scheduler 立刻切换去执行另一个 Warp,用计算掩盖延迟
Warp 调度示意(时间线):

时钟周期:  t1    t2    t3    t4    t5    t6    t7    t8
Warp 0:     计算   计算   等待内存  -     -     计算   计算
Warp 1:     -     计算   计算   计算   等待内存  -     计算
Warp 2:     -     -     -     计算   计算   计算   等待内存
Warp 3:     计算   -     计算   -     计算   -     计算

→ 虽然每个 Warp 都在等内存,但 SM 通过快速切换,让计算单元几乎不停!

3.2.4 寄存器文件(Register File)

  • 每个 SM 有 256 KB 的寄存器文件
  • 分配给该 SM 上运行的所有线程使用
  • 每个线程能分到的寄存器数量有限(如每个线程用太多寄存器,同时运行的线程数就会减少)
  • 延迟极低(~1-2 个时钟周期),是线程的"私人储物柜"

3.2.5 L1 Cache / Shared Memory(共享内存)

这是同一块物理 SRAM,可以通过配置划分给 L1 Cache(硬件自动管理)或 Shared Memory(程序员手动控制)使用。

模式谁管理用途
L1 Cache硬件自动缓存频繁访问的全局内存数据
Shared Memory程序员手动(__shared__同 Block 内线程高速共享数据
  • H100 每个 SM 总共 256 KB,可配置比例
  • 延迟约 20-30 周期,比 HBM 快 20~40 倍
  • 编写高性能 CUDA Kernel 的关键就是善用 Shared Memory

3.3 SM 数量对比:越多越好吗?

GPU架构SM 数量FP32 Core/SMTensor Core/SM总 FP32 Core
V100Volta806485,120
A100Ampere1086446,912
H100Hopper132128416,896
RTX 4090Ada Lovelace128128416,384

SM 数量只是决定性能的因素之一。SM 内部的效率(如 Tensor Core 代数、寄存器带宽、缓存大小)同样重要


四、CUDA Core 与 Tensor Core 的深度协作

4.1 分工图:一个矩阵乘法谁在干活?

以混合精度 GEMM(D = A × B + C,A/B 是 FP16,C/D 是 FP32)为例:

CPU 侧(Host):
  1. 在 CPU 内存准备 A、B、C 数据
  2. cudaMemcpy H2D:把数据搬到 GPU 显存(HBM)
  3. 启动 Kernel <<<grid, block>>>

GPU 侧(Device):
  ┌──────────────────────────────────────────────────────────────┐
  │  Step 1: 数据搬运(Load/Store 单元 + CUDA Core 计算地址)     │
  │  HBM → L2 Cache → L1 Cache/Shared Memory → 寄存器            │
  │                                                              │
  │  Step 2: 矩阵乘法(Tensor Core 专业处理)                     │
  │  寄存器中的小块矩阵 → Tensor Core MMA 运算 → 结果写回寄存器    │
  │                                                              │
  │  Step 3: 累加与写回(CUDA Core 辅助 + Load/Store)            │
  │  FP32 累加 → 结果从寄存器 → Shared Memory → L2 → HBM        │
  └──────────────────────────────────────────────────────────────┘
  
CPU 侧:
  4. cudaMemcpy D2H:把结果 D 搬回 CPU

4.2 算力对比:Tensor Core 为什么快这么多?

以 H100 SXM 为例:

运算单元FP16 算力(Dense)相对速度
CUDA Core(所有 FP32 Core 一起)~67 TFLOPS1×(基准)
Tensor Core(第四代)989 TFLOPS~15×

这意味着:如果一个深度学习算子能用 Tensor Core 跑,就绝对不要用 CUDA Core 跑。

4.3 什么时候 CUDA Core 不可替代?

场景为什么不能用 Tensor Core
数据预处理/后处理不是矩阵运算
索引计算、地址偏移整数运算、逻辑运算
if/else 分支逻辑Tensor Core 只做矩阵乘加
LayerNorm/Softmax 中的归约求和不是矩阵乘法
自定义非线性激活函数Element-wise 运算

五、主流 GPU 规格深度对比:A100 / H100 / H200

5.1 核心规格总表

规格参数A100 (Ampere)H100 (Hopper)H200 (Hopper)
发布时间2020 年 5 月2022 年 3 月2024 年 Q2
架构代号AmpereHopperHopper(同 H100)
制程工艺TSMC 7nmTSMC 4N(5nm 优化版)TSMC 4N
晶体管数量542 亿800 亿800 亿
SM 数量108132132
CUDA Core(FP32)总数~6,912~16,896~16,896
Tensor Core第三代第四代第四代
FP64(双精度浮点)9.7 TFLOPS34 TFLOPS34 TFLOPS
FP32(单精度浮点)19.5 TFLOPS67 TFLOPS67 TFLOPS
TF32 Tensor Core156 TFLOPS494 TFLOPS494 TFLOPS
FP16/BF16 Tensor Core312 TFLOPS / 624*989 TFLOPS / 1,979*989 TFLOPS / 1,979*
FP8 Tensor Core不支持1,979 TFLOPS1,979 TFLOPS
显存类型HBM2eHBM3HBM3e
显存容量40 GB / 80 GB80 GB141 GB
显存带宽1.6 TB/s / 2.0 TB/s3.35 TB/s4.8 TB/s
L2 Cache40 MB50 MB50 MB
TDP(功耗)400W700W700W
NVLink 带宽600 GB/s900 GB/s900 GB/s
特色技术MIG、TF32、结构化稀疏性Transformer Engine、FP8、DPX、TMA同 H100,显存大升级

* 带稀疏性加速的峰值算力。实际应用中通常用 Dense(非稀疏)数值。

5.2 关键解读:H200 为什么只升级显存?

H200 与 H100 算力完全相同,但显存容量从 80GB 提升到 141GB,带宽从 3.35 TB/s 提升到 4.8 TB/s

这传递了一个明确的信号:在大模型时代,显存容量和带宽已经成为比算力更稀缺的资源。

NVIDIA 官方数据:

  • GPT-3 推理:H200 比 H100 快 1.6 倍
  • GPT-3 训练:H200 比 H100 快 ~10%

这些提升完全来自显存升级,而非算力提升。这印证了 Memory Wall 的严重性。

5.3 数据精度科普:FP32、TF32、FP16、BF16、FP8 是什么?

初学者看到这么多精度容易晕,这里用一张表说清楚:

FP32(单精度):  符号位1 + 指数位8 + 尾数位23 = 32位
                  范围大,精度高,但速度慢、占显存
                  
TF32(TensorFloat):符号位1 + 指数位8 + 尾数位10 = 19位有效
                  A100 引入,指数范围和 FP32 一样(不会溢出)
                  尾数减少但用 Tensor Core 加速,训练常用
                  
FP16(半精度):  符号位1 + 指数位5 + 尾数位10 = 16位
                  速度快一倍,但动态范围小,容易梯度下溢
                  
BF16(Brain Float):符号位1 + 指数位8 + 尾数位7 = 16位
                  Google 提出,指数范围和 FP32 一样
                  牺牲精度换稳定性,Transformer 训练友好
                  
FP8(8位浮点):  符号位1 + 指数位4/5 + 尾数位3/2 = 8位
                  H100 引入,推理和部分训练场景
                  配合 Transformer Engine 动态切换精度
精度每参数占用相对 FP32 速度适用场景
FP324 字节通用计算、对精度敏感的场景
TF324 字节(内部 19 位)8×(Tensor Core)训练默认
BF162 字节16×(Tensor Core)训练主力(LLM)
FP162 字节16×(Tensor Core)训练(需 Loss Scaling)
FP81 字节32×(Tensor Core)推理主力、部分训练

六、Memory Wall:为什么显存带宽往往比算力更致命?

6.1 核心矛盾:算力暴涨,带宽瘸腿

过去十年,GPU 峰值算力增长了约 1000 倍,但显存带宽只增长了约 10 倍。两者的鸿沟越来越大,形成了计算机体系结构领域著名的 “Memory Wall”(内存墙)

指标~2012 (GTX 680)~2020 (A100)~2022 (H100)增长倍数
FP32 算力3 TFLOPS19.5 TFLOPS67 TFLOPS~22×
显存带宽192 GB/s2,000 GB/s3,350 GB/s~17×
算力/带宽比~16~10~20基本持平

虽然绝对带宽在增长,但算力增长更快,导致单位算力能分到的带宽反而在减少。

6.2 通俗类比:超级厨师与慢速传送带

想象一家顶级餐厅的后厨:

  • GPU 的算力 = 厨师的刀工和火候。现在这些厨师都是米其林三星水平,一个人能同时处理十几口锅,出菜速度极快。
  • GPU 的显存带宽 = 食材传送带的速度。负责把冰箱里的食材运到厨师手边。

现状是:厨师的刀工越来越快,但食材传送带的速度几乎没变。结果是——

厨师 80% 的时间不是在做菜,而是站在原地等菜上桌。

在 GPU 里,这个"等待"就是:Tensor Core / CUDA Core 处于空闲状态,因为数据还没从 HBM 送达寄存器。无论你的算法多高效、模型多精简,只要数据供应不上,算力就是纯纯的浪费。

6.3 Roofline 模型:判断你的程序卡在哪个瓶颈

Roofline 模型是分析程序性能瓶颈的经典工具。

定义两个概念:

  1. 算术强度(Arithmetic Intensity) = 计算量(FLOP) / 访存量(Byte)

    • 单位:FLOPs/Byte(每读 1 字节数据,做多少次浮点运算)
  2. Machine Balance = 峰值算力 / 峰值带宽

    • 表示硬件的"理想计算/访存比"

以 H100 为例:

FP16 Tensor Core 峰值算力:989 TFLOPS = 989,000 GFLOPS
HBM 峰值带宽:3.35 TB/s = 3,350 GB/s

Machine Balance = 989,000 / 3,350 ≈ 295 FLOPs/Byte

含义:H100 上,只有当每读 1 字节数据做 295 次以上浮点运算时,算力才能跑满。低于这个值,瓶颈就在显存带宽

Roofline 图(概念版):

性能 (GFLOPS)
989 ├─────────────────────────────┐ ← 算力上限(平顶)
    │                             │
500 ├──────────────┐              │
    │            / │              │
200 ├──────────/───┘              │
    │        /  ← 带宽受限斜线(斜率 = 带宽)
100 ├──────/                      │
    │    /                        │
 50 ├──/                          │
    │/                            │
 10 ├─────────────────────────────┘
    └────┬────┬────┬────┬────┬────→ 算术强度 (FLOPs/Byte)
         1    10   50  100  295
         Memory Bound 区域      Compute Bound 区域
         (斜线部分)            (平顶部分)

常见算子的算术强度:

算子算术强度瓶颈类型说明
Element-wise(如 ReLU)~0.1Memory Bound读一次算一次,几乎不重用数据
LayerNorm / Softmax~1-5Memory Bound需要多次读取同一数据进行归约
Attention(标准实现)~5-20Memory Bound频繁读写巨大的 $N \times N$ 矩阵
GEMM(大矩阵乘法)>100Compute Bound数据在寄存器/Shared Memory 中被大量复用

这就是为什么 FlashAttention 能带来数倍加速:它没有提升算力,而是通过分块计算把 Attention 从 Memory Bound 往 Compute Bound 方向推。

6.4 大模型时代,Memory Wall 更严重了

原因详细说明
模型参数爆炸GPT-4 级别模型参数量达万亿级,仅权重就需要数 TB 存储。单卡 80GB 根本放不下。
激活值巨大长序列(32K、128K、甚至 1M tokens)导致中间激活值占用海量显存。
Batch Size 受限显存不够大,无法通过增大 Batch 来提升数据复用率,导致有效带宽利用率低。
推理阶段尤为明显推理是逐 token 生成(自回归),每次只算一个位置,计算/访存比极低,严重 Memory Bound。
通信瓶颈多卡训练时,模型并行/数据并行需要在卡间传输数据,进一步加剧带宽压力。

结论:在 A100/H100/H200 上,显存带宽瓶颈往往比算力瓶颈更致命。这也是为什么 FlashAttention、PagedAttention、KV Cache 优化、量化(INT8/FP8/INT4)等减少访存或提升有效带宽的技术如此重要。


七、GPU 存储层次结构:写 Kernel 的核心艺术

GPU 的存储系统是一个典型的金字塔结构:越靠近顶端的存储,速度越快、容量越小、越珍贵;越靠近底端,速度越慢、容量越大。

7.1 完整层次金字塔

        速度 ↑
    极快     │      ┌─────────────┐
             │      │   寄存器     │  ← 线程私有,~1-2 周期,~256 KB/SM
    快       │      ├─────────────┤
             │      │ L1 / Shared │  ← Block 共享,~20-30 周期,~128-256 KB/SM
    中等     │      ├─────────────┤
             │      │   L2 Cache  │  ← GPU 共享,~100-300 周期,~40-50 MB
    慢       │      ├─────────────┤
             │      │     HBM     │  ← 全局显存,~400-800 周期,80-141 GB
    极慢     │      ├─────────────┤
             │      │  主机内存    │  ← 经 PCIe,数千周期,数百 GB
             │      └─────────────┘
             └────────────────────────→ 容量

7.2 各层存储详解

7.2.1 寄存器(Registers)—— 最快的私人储物柜

属性详情
归属每个线程私有
物理位置SM 内部,紧邻 ALU(计算单元)
典型延迟1-2 个时钟周期(几乎瞬时)
容量H100 每 SM 256 KB,由该 SM 上所有线程共享分配
编程方式编译器自动分配,程序员通过局部变量间接使用

类比:寄存器就像你口袋里的小笔记本——触手可及,但容量极小。

关键限制

  • 每个线程能用的寄存器是有限的(如 H100 最多 255 个 32 位寄存器/线程)
  • 如果用的寄存器太多,SM 上能同时驻留的线程数就会减少(Occupancy 降低)
  • 如果寄存器不够用,会发生 Register Spill,数据被挤出到 Local Memory(本质上是 HBM),性能暴跌

7.2.2 L1 Cache / Shared Memory(共享内存)—— 车间的公共白板

属性详情
归属同一线程块(Block)内的线程共享
物理位置SM 内部,片上 SRAM
典型延迟20-30 个时钟周期(比 HBM 快 20-40 倍)
容量H100 每 SM 256 KB(L1 + Shared 合计,可配置比例)
编程方式L1 由硬件自动管理;Shared Memory 由程序员显式控制(__shared__

类比:Shared Memory 就像车间里的公共白板——同一个车间的工人都能快速读写,但不同车间之间看不到彼此的白板。

为什么 Shared Memory 是优化关键?

场景:矩阵乘法 C = A × B

Naive 做法(只用 HBM):
  每个线程读取 A 的一行和 B 的一列,计算 C 的一个元素。
  A 的每一行被读取 N 次,B 的每一列被读取 M 次。
  → 大量重复访存,HBM 带宽成为瓶颈。

Tiled 做法(使用 Shared Memory):
  1. 将 A 和 B 分成小块(Tile)。
  2. 一个 Block 的线程协作,把 A 的一个 Tile 和 B 的一个 Tile 从 HBM 加载到 Shared Memory。
  3. 每个线程从 Shared Memory(快)读取数据计算,而不是从 HBM(慢)。
  4. 算完当前 Tile,再加载下一个 Tile。
  → A/B 的每个元素被复用多次,HBM 访存次数减少 BlockSize 倍!

Shared Memory 的 Bank Conflict(银行冲突):

Shared Memory 被分成 32 个 Bank(类似 32 条并行通道)。理想情况下,一个 Warp(32 线程)同时访问 32 个不同 Bank,可以一次完成。但如果多个线程同时访问同一个 Bank,就会发生冲突,访问被串行化,性能下降。

理想情况(无 Bank Conflict):
  线程 0 读 Bank 0,线程 1 读 Bank 1,...,线程 31 读 Bank 31
  → 一次访问全部完成!

冲突情况(Bank Conflict):
  线程 0 和线程 1 同时读 Bank 0
  → 需要两次访问,性能减半

7.2.3 L2 Cache —— 全厂的中转仓库

属性详情
归属整个 GPU 所有 SM 共享
物理位置GPU 芯片上,但不在 SM 内部
典型延迟100-300 个时钟周期
容量A100: 40 MB,H100: 50 MB,B200: 126 MB
编程方式硬件自动管理,程序员无法直接控制

作用

  • 缓存 HBM 数据,减少 SM 对 HBM 的直接访问
  • 承担 SM 之间的数据交换(一个 SM 计算的结果被另一个 SM 读取时)
  • 原子操作(Atomic)的同步点之一

L2 Cache 从 H100 到 B200 暴涨到 126 MB,这对小 Batch 推理特别重要——因为小 Batch 时激活值较小,更容易全部留在 L2 里,避免访问 HBM。

7.2.4 HBM(High Bandwidth Memory)—— 大仓库,但路远

属性详情
归属全局显存,所有线程可访问
物理位置GPU 芯片外,通过硅中介层(Interposer)2.5D/3D 封装连接
典型延迟400-800 个时钟周期(是寄存器的 400 倍
容量A100: 80 GB,H100: 80 GB,H200: 141 GB
带宽A100: 2.0 TB/s,H100: 3.35 TB/s,H200: 4.8 TB/s
编程方式cudaMalloc 分配的内存就是 HBM;Kernel 中直接读写全局变量

为什么 HBM 这么快还叫"慢"?

HBM 的"快"是相对于传统 DDR/GDDR 而言的。但和片上存储(寄存器、Shared Memory)相比,HBM 仍然慢了 2~3 个数量级。在 GPU 的世界里,延迟 400 周期就是"慢"

HBM 访问优化——合并访问(Coalesced Access):

Warp 里的 32 个线程如果同时访问连续的内存地址,HBM 控制器可以把这 32 个请求合并成一次宽事务(如 128 字节),效率极高。

合并访问(高效):
  线程 0 读 address[0],线程 1 读 address[1],...,线程 31 读 address[31]
  → 合并为一次 128 字节事务,带宽利用率 100%

非合并访问(低效):
  线程 0 读 address[0],线程 1 读 address[1000],线程 2 读 address[2000]...
  → 32 次独立事务,带宽利用率暴跌

7.2.5 主机内存(Host Memory)—— 另一个城市的仓库

属性详情
归属CPU 主存(DRAM)
连接方式通过 PCIe 总线与 GPU 连接
典型延迟数千到数万时钟周期(需经过 PCIe + 驱动 + 系统调用)
带宽PCIe 4.0 x16: ~32 GB/s,PCIe 5.0 x16: ~64 GB/s
容量数百 GB 到数 TB

PCIe 带宽只有 HBM 的 1/50 ~ 1/100。因此:尽量减少 CPU↔GPU 的数据传输。大模型训练时,模型权重和激活值应尽量常驻 GPU 显存,不要频繁来回搬运。

7.3 各层存储对比速查表

存储层级延迟(周期)相对 HBM 速度容量/单元管理方
寄存器1-2~400×~256 KB/SM编译器自动
L1/Shared Memory20-30~20×~128-256 KB/SM硬件 / 程序员
L2 Cache100-300~3-5×~40-50 MB/GPU硬件自动
HBM400-800~80-141 GB/GPU程序员
主机内存数千+~0.02×~数百 GB操作系统

7.4 Flash Attention:存储层次优化的典范

Flash Attention 是理解 Memory Wall 和存储层次优化的最佳案例:

标准 Attention 的问题:

输入: Q, K, V(各为 N×d 矩阵)
步骤 1: S = Q × K^T      → 产生 N×N 中间矩阵,写回 HBM
步骤 2: P = softmax(S)   → 读取 S,计算 P,写回 HBM
步骤 3: O = P × V        → 读取 P 和 V,计算 O

问题:N×N 矩阵(如 N=32768, 那 N×N = 1B 个元素,FP16 占 2GB)
      在 HBM 中来回读写 3-4 次,严重 Memory Bound!

Flash Attention 的解法:

核心思想:分块计算 + 增量 Softmax

1. 将 Q, K, V 切成小块(Tile),加载到快速的 SRAM(Shared Memory + 寄存器)。
2. 在 SRAM 内完成:
   - 计算当前块的 S_tile = Q_tile × K_tile^T
   - 增量更新 softmax 统计量(Online Softmax),不保存完整 S 矩阵
   - 计算当前块的输出 O_tile
3. 只把最终输出 O 写回 HBM,中间的 S 和 P 根本不存!

效果:HBM 访存量从 O(N^2) 降到 O(N),在带宽受限场景下获得数倍加速。

这完美印证了本文的核心观点:写高性能 Kernel,就是在这个存储层次结构中最大化数据复用,让数据尽量在快速层级(寄存器、Shared Memory)中"多待一会儿",减少对慢速 HBM 的访问


八、编程实践:如何根据硬件特性写高效 Kernel?

8.1 Occupancy(占用率):你的 SM 有多"满"?

Occupancy = 当前活跃的 Warp 数 / SM 支持的最大 Warp 数

  • H100 每个 SM 最多支持 64 个 Warp(= 2048 线程)
  • 如果 Occupancy 低,意味着 SM 上有很多"空位",当某些 Warp 在等数据时,没有足够多其他 Warp 可切换来掩盖延迟

影响 Occupancy 的因素:

因素影响
每个线程的寄存器用量用的寄存器越多,同时驻留的线程越少
Shared Memory 用量一个 Block 用的 Shared Memory 越多,同时驻留的 Block 越少
Block 大小Block 太小(如 32)可能导致 Warp 数不足
每个 Block 的线程数需是 Warp 大小(32)的整数倍

8.2 高效 Kernel 检查清单

□ 1. 内存合并访问
    → Warp 内线程访问连续地址,最大化 HBM 带宽利用率

□ 2. 避免 Bank Conflict
    → Shared Memory 访问模式让 32 个线程命中不同 Bank

□ 3. 使用 Shared Memory 做数据复用
    → 一次从 HBM 加载,多次在 Shared Memory 复用(Tiling)

□ 4. 减少寄存器用量
    → 避免 Register Spill 到 HBM

□ 5. 控制分支发散
    → 一个 Warp 内避免 if/else 导致线程走不同路径

□ 6. 足够的 Occupancy
    → 让每个 SM 上有足够多的 Warp 可供调度,掩盖访存延迟

□ 7. 计算与通信重叠
    → 使用 CUDA Stream 和异步拷贝,让计算和访存并行

九、初学者常见误区

误区真相
“GPU 线程数越多越好”线程太多不一定更快,还要看内存访问模式和数据复用。而且一个 SM 同时执行的 Warp 数有上限。
“Tensor Core 万能”Tensor Core 只做矩阵乘加。LayerNorm、Softmax、Embedding 等非矩阵运算它不管。
“HBM 带宽够用了”实际有效带宽通常只有峰值的 60-80%。如果访问模式不好(非合并),可能只有 10-20%。
“L1/L2 Cache 会自动帮我优化”Cache 对规则访问模式有效,但对随机访问(如 Attention 中的稀疏索引)帮助有限,仍需手动用 Shared Memory。
“FP16 一定比 FP32 快”只有 Tensor Core 做矩阵乘法时 FP16 才快。如果代码跑在 CUDA Core 上,FP16 未必更快,甚至可能因为类型转换更慢。
“买更多 GPU 就能线性加速”多卡之间有通信瓶颈(NVLink/网络),模型并行/数据并行都会引入通信开销,扩展效率 rarely 100%。

十、本章知识自测

试着回答以下问题,检验理解程度:

  1. SM 里有哪些核心组件?CUDA Core 和 Tensor Core 各自负责什么?
  2. 为什么 H200 算力和 H100 一样,但推理能快 1.6 倍?
  3. 用"厨师与传送带"类比,解释什么是 Memory Wall。
  4. GPU 存储层次从快到慢依次是什么?相邻两层速度大约差多少倍?
  5. 什么是合并访问(Coalesced Access)?为什么它很重要?
  6. Roofline 模型中,Arithmetic Intensity 怎么算?H100 的 Machine Balance 约是多少?
  7. Flash Attention 减少了哪类存储的访问?它是如何做到的?
  8. 为什么写 CUDA Kernel 时要关注 Occupancy?

十一、总结与核心口诀

维度核心要点
SMGPU 的计算心脏,包含 CUDA Core(通用)+ Tensor Core(专用矩阵)+ 寄存器 + Shared Memory + Warp Scheduler
Tensor CoreAI 算力的主要来源,专精 GEMM,比 CUDA Core 快约 15 倍。从 Volta 到 Blackwell 持续演进,支持精度越来越低(FP16→BF16→FP8→FP4)
CUDA Core通用计算单元,负责标量/向量运算、地址计算、逻辑控制,不可替代
Memory Wall算力增长(1000×)远超带宽增长(10×),计算单元大部分时间空等数据,是大模型时代的主要瓶颈
存储层次寄存器 > Shared Memory > L2 Cache > HBM > 主机内存。延迟相差可达 1000 倍。Kernel 优化的本质是最大化片上数据复用
带宽瓶颈H200 只升级显存就获得 1.6× 推理提升,证明带宽比算力更稀缺

🎯 黄金法则

“能留在片上(寄存器 / Shared Memory)的数据,绝不回 HBM;能合并的访存,绝不分散;能被多个线程复用的数据,只从 HBM 加载一次。”

理解这些硬件原理,是后续学习 CUDA 编程、算子优化(如 Triton、CUTLASS)、分布式训练(Tensor/Pipeline Parallelism)和推理加速(vLLM、TensorRT-LLM)的绝对基石


参考与延伸阅读

评论区