📑 本页目录（点开跳转）

🎯 一句话
🆚 一、CPU vs GPU
🏗️ 二、GPU 的层次结构
🧵 三、执行模型：warp 和它的两个坑
📊 四、几张主流卡的关键数字
🧰 五、你实际会碰到的抽象层
🔗 和站内其他章的关系
✅ 检查点
🛑 可以停在这里

02 · GPU 到底是什么

⏱ 26 分钟 ｜ ⭐ 不需要写 CUDA，但要知道它长什么样

🎯 一句话

CPU 是几个博士生，GPU 是几千个小学生。 问题不在于谁更聪明，而在于你的任务能不能拆成几千份一模一样的小活。 深度学习恰好完美符合 —— 这就是 GPU 统治这个领域的全部原因。

🆚 一、CPU vs GPU

CPU： GPU：

⭐ 省下来的晶体管去哪了：GPU 不做分支预测、不做乱序执行，把面积全换成了核心和寄存器 —— 代价是它靠切换线程掩盖内存延迟，而不是靠缓存。

	CPU	GPU
核心数	8 ~ 128	几千个
单核性能	强	弱
设计目标	降低延迟（把一件事尽快做完）	提高吞吐（同时做很多事）⭐
擅长	分支多、依赖强的逻辑	大量相同的、独立的运算
缓存	大（几十 MB）	小（每 SM 几百 KB）

🔑 一个关键的设计取舍： GPU 把省下来的晶体管（不做分支预测、不做乱序执行）全用来堆核心和寄存器。 代价是：它靠"切换到别的线程"来掩盖内存延迟，而不是靠缓存。

🏗️ 二、GPU 的层次结构

先看包含关系：运算单元与局部存储在 SM 内；再看共享关系：多个 SM 连接 L2，继续访问显存。图示是逻辑层次，不是芯片物理布局。

图下说明

GPU
SM（流多处理器）× 100+ 真正干活的单元
CUDA Core （FP32 通用运算）
Tensor Core （⭐ 矩阵乘专用，深度学习的主力）
寄存器（最快，每 SM 几百 KB）
Shared Memory / L1（几百 KB，⭐ FlashAttention 靠它）
L2 Cache（几十 MB，全卡共享）
HBM 显存（40 / 80 / 141 GB）慢，但容量大

⭐ Tensor Core：为什么它改变了一切

结果对照

CUDA Core：一次做一个乘加 a×b+c

Tensor Core：一次做一整个【小矩阵乘】 D = A×B + C（比如 4×4）

单位时间的矩阵乘吞吐高【一个数量级】

A100 的对照：

FP32（CUDA Core） 19.5 TFLOPS

TF32（Tensor Core） 156 TFLOPS ⭐ 8 倍

BF16/FP16（Tensor Core） 312 TFLOPS ⭐ 16 倍

INT8（Tensor Core） 624 TOPS

🔑 这解释了第 6 章混合精度为什么是"免费的午餐"： 不是因为半精度算得快，而是因为半精度才能用上 Tensor Core。

⚠️ 但 Tensor Core 有使用条件：矩阵维度要对齐（通常是 8 或 16 的倍数）。 维度是 4095 而不是 4096，可能直接掉回 CUDA Core —— 这是个真实的坑。

🧵 三、执行模型：warp 和它的两个坑

从外框向内读组织层级；线程格仅列出首尾示例，不是说一个 warp 只有三个线程。这张图不表示执行先后。

图下说明

线程（thread）
32 个线程打包成一个【warp】 ⭐ 这是真正的调度单位
若干 warp 组成一个 block（跑在同一个 SM 上）
若干 block 组成 grid
⭐ 一个 warp 里的 32 个线程【必须执行同一条指令】（SIMT）

坑 ① warp divergence（分支发散）

结果对照

if (threadIdx.x % 2 == 0) A();

else B();

同一个 warp 里既有走 A 的又有走 B 的

硬件只能【先跑 A（走 B 的线程闲置），再跑 B】

耗时翻倍 💀

⭐ 这就是为什么 GPU 不擅长分支多的逻辑

坑 ② 非合并访存（uncoalesced access）

结果对照

✅ 合并访存：warp 里 32 个线程访问【连续】的地址

硬件合并成 1 次内存事务

❌ 非合并：访问跳跃的地址（比如按列访问行主序矩阵）

可能变成 32 次独立事务→有效带宽掉到 1/32 💀

⭐ 这就是为什么【张量的内存布局（contiguous）很重要】

# PyTorch 里的真实影响
import torch
x = torch.randn(4096, 4096, device='cuda')
y = x.t()                      # 转置只改 stride，没搬数据
z = y @ w                      # ⚠️ 非连续布局，可能触发隐式拷贝或慢路径

y = x.t().contiguous()         # ⭐ 显式重排一次，后续访问都是合并的

💡 为什么"切换线程"能掩盖延迟：当一个 warp 在等 HBM 数据（几百个周期）时，SM 立刻切换到另一个就绪的 warp。 只要同时驻留的 warp 足够多（occupancy 高），延迟就被填满了。 —— 这也是为什么 batch 太小时 GPU 效率低：warp 不够多，掩盖不住延迟。

📊 四、几张主流卡的关键数字

	A100 80G	H100 SXM	RTX 4090
BF16 算力	312 TFLOPS	989 TFLOPS	165 TFLOPS
FP8 算力	❌	1979 TFLOPS ⭐	❌
显存	80 GB	80 GB	24 GB ⚠️
显存带宽	2.0 TB/s	3.35 TB/s	1.0 TB/s
卡间互联	NVLink 600 GB/s	NVLink 900 GB/s	⚠️ 仅 PCIe ~64GB/s

🔑 看最后两行，不要只看算力： 4090 的算力有 A100 一半，但显存只有 24GB，卡间互联只有 PCIe。 → 单卡跑小模型很划算，多卡训练大模型基本不可用（第 12 章会解释为什么张量并行离不开 NVLink）。

💡 H100 的 FP8 值得注意：算力再翻一倍。但 FP8 的动态范围极窄，需要 per-tensor scaling 才能训得稳（第 6 章）。

🧰 五、你实际会碰到的抽象层

信息关系

你写的： PyTorch （torch.matmul）

框架调用的： cuBLAS / cuDNN / CUTLASS ← NVIDIA 的高度优化库

或者：自定义 kernel（CUDA / Triton）← 第 7 章

驱动→硬件

⭐ 实践建议： 99% 的情况下你不需要写 CUDA。 先确认瓶颈真的在某个算子上（第 4 章教你怎么测），再考虑用 Triton（Python 语法，门槛低得多）—— 第 7 章。

# 三个你一定会用到的命令
nvidia-smi                    # 看显存占用、利用率、温度、功耗
nvidia-smi topo -m            # ⭐ 看卡间互联拓扑（NVLink 还是 PCIe）
nvidia-smi dmon -s pucm       # 持续监控

⚠️ nvidia-smi 的 "GPU-Util" 极具误导性：它只表示"过去一段时间内有 kernel 在跑"的百分比， 不代表算力被用满。一个只用了 5% 算力的 kernel 一直跑，它也显示 100%。 真正该看的是 MFU（第 4 章）。

🔗 和站内其他章的关系

相关的地方	这里的位置
ML 基础第 15 章 `.to('cuda')`	现在你知道那边有什么了
第 1 章算力涨得比带宽快	HBM 就是那个瓶颈
第 3 章	详细讲 HBM / L2 / Shared Memory 的层次
第 6 章混合精度	为了用上 Tensor Core ⭐
第 8 章 FlashAttention	靠 Shared Memory 避免往返 HBM ⭐

✅ 检查点

CPU 和 GPU 的设计目标分别是什么？
GPU 靠什么掩盖内存延迟？这对 batch size 有什么含义？
Tensor Core 和 CUDA Core 的区别？A100 上差多少倍？
混合精度为什么快？（真正的原因）
什么是 warp？warp divergence 为什么慢？
什么是非合并访存？它和 PyTorch 里的什么操作有关？
4090 为什么不适合多卡训练大模型？
为什么 nvidia-smi 的 GPU-Util 具有误导性？

👀 答案

CPU 降低延迟（把一件事尽快做完），GPU 提高吞吐（同时做很多事）。GPU 把不做分支预测/乱序执行省下的晶体管全用来堆核心。
靠切换到别的 warp——一个 warp 等 HBM 时立刻切到另一个就绪的。含义：batch 太小时驻留 warp 不够多，掩盖不住延迟，GPU 效率低。
CUDA Core 一次做一个乘加，Tensor Core 一次做一整个小矩阵乘。A100 上 FP32 是 19.5 TFLOPS，BF16 Tensor Core 是 312 TFLOPS，16 倍。
不是因为半精度算得快，而是因为半精度才能用上 Tensor Core。⚠️ 且要求矩阵维度对齐（8 或 16 的倍数），4095 而不是 4096 可能掉回 CUDA Core。
32 个线程打包成的调度单位，一个 warp 里所有线程必须执行同一条指令（SIMT）。divergence 时硬件只能先跑一个分支（另一批线程闲置）再跑另一个，耗时翻倍。
warp 里的线程访问跳跃的地址（如按列访问行主序矩阵），无法合并成一次内存事务，有效带宽可能掉到 1/32。和 PyTorch 里张量是否 contiguous 直接相关——转置只改 stride 不搬数据，后续访问可能走慢路径。
因为显存只有 24GB，且卡间互联只有 PCIe（~64GB/s）没有 NVLink。张量并行需要极高的卡间带宽，PCIe 撑不住。
它只表示"过去一段时间内有 kernel 在跑"的百分比，不代表算力被用满。一个只用 5% 算力的 kernel 一直跑也显示 100%。真正该看 MFU。

🛑 可以停在这里

⚡ 走神救援

先记住这几件事

GPU 通过大量并行任务提高吞吐；SM 内的执行单元与存储层次需要一起理解。
Tensor Core 面向矩阵运算，实际收益还取决于精度、形状与数据布局。
分支分歧和跳跃访存都会降低效率；检查真实执行与搬运，不只看利用率百分比。

下一节 👉 03-显存与带宽墙.md ⭐

打卡记录保存在你的浏览器里，首页能看到总进度