02 · GPU 到底是什么

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

🎯 一句话

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

🆚 一、CPU vs GPU

CPU： GPU：

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

🏗️ 二、GPU 的层次结构

   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 和它的两个坑

   线程（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 不够多，掩盖不住延迟。

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

🔑 看最后两行，不要只看算力： 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 章）。

🔗 和站内其他章的关系

✅ 检查点

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

🛑 可以停在这里

⚡ 走神救援

CPU 是几个博士生（降低延迟），GPU 是几千个小学生（提高吞吐）；GPU 把不做分支预测/乱序执行省下的晶体管全堆成核心，靠切换 warp 而不是靠缓存来掩盖内存延迟 → ⭐batch 太小时驻留 warp 不够，掩盖不住延迟，效率低。层次：SM（含 CUDA Core / Tensor Core / 寄存器 / Shared Memory 几百KB）→ L2（几十MB）→ HBM（80GB，慢）。⭐Tensor Core 一次做一整个小矩阵乘，A100 上 FP32 19.5 TFLOPS vs BF16 312 TFLOPS（16倍） → ⭐混合精度快的真正原因不是"半精度算得快"，而是半精度才能用上 Tensor Core；⚠️要求维度对齐（4095 而非 4096 可能掉回 CUDA Core）。执行模型：32 线程 = 1 warp，必须执行同一条指令（SIMT）；两个坑：⚠️warp divergence（同 warp 走不同分支 → 先跑一个再跑另一个，耗时翻倍，所以 GPU 不擅长分支）、⚠️非合并访存（访问跳跃地址 → 有效带宽掉到 1/32，所以 PyTorch 里 contiguous 很重要）。关键数字：A100 312 TFLOPS / 2.0TB/s / NVLink 600GB/s；H100 989 TFLOPS（FP8 翻倍到 1979）/ 3.35TB/s / NVLink 900GB/s；⭐4090 算力不差但显存仅 24GB 且只有 PCIe 无 NVLink → 多卡训练大模型基本不可用。⚠️⭐nvidia-smi 的 GPU-Util 极具误导性——它只说"有 kernel 在跑"，一个只用 5% 算力的 kernel 也显示 100%，真正该看 MFU。99% 的情况不需要写 CUDA，真要写就用 Triton。

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