6. GPU与CPU的架构分野:并行计算的硬件哲学
训练同一个神经网络,GPU 比 CPU 快 100 倍,为什么?
答案要从 CPU 的困境说起
第一章:CPU 的困境
1.1 功耗墙:单核性能的物理天花板
2004 年,Intel 取消了 7GHz 的 Pentium 4 后继项目 Tejas。原因很简单:芯片功耗将超过 150W,即使最先进的散热系统也无法稳定运行。
背后的物理规律:
动态功耗 = α × C × V² × f
其中:
α = 开关活动因子 (0.1-0.3)
C = 负载电容
V = 供电电压
f = 时钟频率
关键问题:
• 从公式看:P ∝ V² × f(线性于频率)
• 但实际中:提升频率必须同时提升电压(否则晶体管无法稳定开关)
• 经验关系:V ∝ f(近似)
• 综合效果:P ∝ f² × f = f³
实例:频率提升2× → 电压提升约1.4× → 功耗增加 2 × 1.4² ≈ 4×
(实际由于漏电流等因素,接近6-8×)
为什么不能通过更好的散热解决?
功耗 150W 不仅是散热问题,还涉及:
| 问题 | 影响 | 数据 |
|---|---|---|
| 电源传输 | 需要数百个电源引脚 | 150W@1.2V = 125A 电流 |
| 压降(IR Drop) | 芯片远端电压不足 | 可能从 1.2V 降到 0.9V → 性能大降 |
| 热密度 | 接近核反应堆水平 | 75-150 W/cm² |
结论:频率竞赛终结,处理器进入多核时代。
1.2 多核的尝试:阿姆达尔定律的阴影
Intel 和 AMD 转向多核,但很快遇到新瓶颈:不是所有代码都能并行。
核心思想(直白版)
假设你的程序有这样的结构:
程序总耗时:100秒
├─ 90秒:可以并行的部分(如矩阵计算)
└─ 10秒:必须串行的部分(如读取配置文件、更新全局状态)
用2核CPU运行:
├─ 并行部分:90秒 / 2 = 45秒
└─ 串行部分:10秒(不变!)
总时间:45 + 10 = 55秒
加速比:100/55 = 1.82×(不是理想的2×!)
用无限多核运行:
├─ 并行部分:90秒 / ∞ = 0秒
└─ 串行部分:10秒(仍然不变!)
总时间:10秒
加速比:100/10 = 10×(理论上限!)
关键洞察:串行代码成为性能的天花板,再多核心也无济于事。
阿姆达尔定律(数学表达)
加速比 = 1 / (S + P/N)
S = 串行代码占比
P = 可并行代码占比 (S+P=1)
N = 核心数
实际应用的串行占比
不同类型的应用,串行占比差异巨大:
| 应用类型 | 串行占比 | 4 核加速 | 16 核加速 | ∞ 核加速 | 瓶颈原因 | 例子 |
|---|---|---|---|---|---|---|
| 通用应用 | 20-50% | 1.6-2.5× | 1.8-3.2× | 2-5× | 大量串行逻辑 | 浏览器、操作系统 |
| 数据库 | 10-20% | 3.1-3.5× | 4.7-6.4× | 5-10× | 事务管理、锁 | MySQL、PostgreSQL |
| 科学计算 | 1-5% | 3.8-3.9× | 13.9-15.2× | 20-100× | 少量初始化代码 | 矩阵运算、模拟 |
| 深度学习 | <1% | 4.0× | 15.8× | >100× | 几乎纯并行 | 神经网络训练 |
关键数据:
- 桌面应用(浏览器、Office):串行占比 30-50% → 多核收益小
- 服务器应用(数据库):串行占比 10-20% → 16 核后收益递减
- 深度学习:串行占比<1% → CPU 的多核策略不够用
这就是为什么消费级 CPU 很少超过 16 核——对于大多数应用,增加更多核心的边际收益递减。
1.3 CPU 的死胡同
让我们总结 CPU 在 2000-2010 年代面临的困境:
性能提升路径
│
├─ 单核频率 ✗ (功耗墙)
│ └─ 已达到物理极限 (~5GHz)
│
└─ 多核并行 △ (阿姆达尔定律)
└─ 受串行代码限制
└─ 大多数应用:10-50%串行
└─ 理论上限:2-10×加速
关键问题:如果应用有 50%串行代码,即使有 1000 个核心,最多也只能加速 2 倍。
但是,有一类应用不受此限制:图形渲染和科学计算。这些 workload 的并行度可达 99%甚至 99.9%。
这就是 GPU 的机会窗口。
第二章:两种哲学——延迟优先 vs 吞吐量优先
在深入技术细节前,先理解 CPU 和 GPU 的核心设计差异:
2.1 设计目标对比
| 维度 | CPU 哲学 | GPU 哲学 | 权衡 |
|---|---|---|---|
| 首要目标 | 最小化单线程延迟 | 最大化总吞吐量 | - |
| 适用场景 | 串行代码、复杂逻辑 | 数据并行、简单逻辑 | - |
| 核心数量 | 少 (8-64 核) | 多 (10,000+核) | 面积/功耗 |
| 单核复杂度 | 高 (乱序执行、预测) | 低 (顺序执行) | 控制开销 |
| 延迟隐藏 | 乱序+预测 | 海量线程切换 | 硬件成本 |
| 编程模型 | 隐式并行 (编译器) | 显式并行 (程序员) | 易用性 |
核心洞察:CPU 和 GPU 不是"谁更好",而是为不同问题优化。
2.2 具体场景:1000 次浮点乘加(FMA)
让我们用一个具体任务对比两者的策略:
任务:计算 result[i] = a[i] * b[i] + c[i],共 1000 次独立计算
CPU 策略 (Intel Core i9-13900K, 24 核心)
单核性能:
├─ 每核2个FMA单元
├─ 乱序执行:同时处理200条指令
└─ 完成1000次FMA:~500周期 = 0.1微秒 @ 5GHz
全核性能 (24核):
└─ 1000次FMA:~21周期 = 4.2纳秒
GPU 策略 (NVIDIA H100, 16,896 个 CUDA Core)
单核性能:
├─ 简单顺序流水线 (4级)
├─ 单个CUDA Core完成1次FMA:4周期
└─ 完成1000次FMA:4000周期 = 2.2微秒 @ 1.8GHz
(比CPU单核慢22×!)
全核性能 (16,896核):
└─ 1000次FMA:ceil(1000/16896) × 4 = 4周期 = 2.2纳秒
(比CPU快2×)
扩大规模到 100 万次 FMA:
| 处理器 | 单核延迟 | 全核时间 | GPU 优势 |
|---|---|---|---|
| CPU | 0.1 微秒 | 4,167 微秒 | - |
| GPU | 2.2 微秒 | 33 微秒 | 126× |
这就是"牺牲延迟换吞吐量"的本质:
- GPU 单线程慢(22×)
- 但并行度高(704×),最终吞吐量高(126×)
2.3 晶体管预算:设计哲学的"投票"
理念再好也需要硬件落地。让我们看看 CPU 和 GPU 如何用晶体管"投票"表达设计哲学。
CPU 的晶体管分配 (Intel Core i9, 单核 8 亿晶体管)
总计:~8亿晶体管/核心
分配:
缓存 (L1+L2) ██████████████████████████████████████ 40% (3.2亿)
└─ 目的:减少访问主存,降低延迟
乱序执行引擎 ███████████████████████ 25% (2.0亿)
└─ 重排序缓冲区(ROB)、寄存器重命名
└─ 目的:隐藏指令依赖,提升IPC
分支预测器 ██████████ 10% (0.8亿)
└─ 全局/局部历史、混合预测器
└─ 目的:避免流水线停顿
计算单元 (ALU/FPU) ██████████ 10% (0.8亿)
└─ 实际执行运算
其他控制逻辑 ███████████████ 15% (1.2亿)
关键观察:只有 10%晶体管用于计算,90%服务于"降低延迟"!
GPU 的晶体管分配 (H100 单个 SM, 约 6 亿晶体管)
总计:~6亿晶体管/SM
分配:
CUDA Core (计算) ████████████████████████████████████████████████████ 60% (3.6亿)
└─ 128个FP32/INT32核心
└─ 目的:最大化计算密度
Tensor Core ███████████████ 15% (0.9亿)
└─ 专用矩阵运算加速器
寄存器堆 ██████████ 10% (0.6亿)
└─ 256KB,存储数千线程状态
Shared Memory/L1 ████████ 8% (0.48亿)
└─ 228KB,显式管理的快速缓存
控制逻辑 (调度器等) █████ 5% (0.3亿)
└─ 简化的warp调度
其他 ██ 2% (0.12亿)
关键差异:60%晶体管用于计算,控制逻辑仅 5%!
直接对比
| 资源类型 | CPU 占比 | GPU 占比 | GPU/CPU 比率 | 含义 |
|---|---|---|---|---|
| 计算单元 | 10% | 60% | 6× | GPU 计算密度高 |
| 控制逻辑 | 50% | 5% | 0.1× | GPU 控制开销低 |
| 缓存 | 40% | 18% | 0.45× | GPU 更依赖程序员管理 |
So What:相同芯片面积下,GPU 能塞入 6 倍的计算单元。这是性能差异的根本原因。
第三章:CPU 的延迟隐藏技术——复杂但有效
理解了设计哲学后,让我们深入 CPU 如何用 50%的晶体管"服务"10%的计算单元。
3.1 问题:内存延迟的挑战
现代 CPU 访问主存的延迟约300 周期 (100ns)。如果 CPU 在等待数据时停顿,性能会暴跌。
int x = array[i]; // 从内存加载,需要300周期
int y = x * 2; // 依赖x,必须等待
CPU 通过三大技术隐藏这 300 周期:
3.2 技术 1:乱序执行 (Out-of-Order Execution)
核心思想:CPU"偷看"后面的指令,找到与当前等待无关的指令提前执行。
原始指令序列:
1. x = array[i] // 加载,需要300周期
2. y = x * 2 // 依赖x,必须等待
3. a = b + c // 与x无关!
4. d = e * f // 与x无关!
乱序执行后:
1. 开始加载 x = array[i] (周期0)
2. 立即执行 a = b + c (周期1,不等x)
3. 立即执行 d = e * f (周期2,不等x)
4. ... (执行更多无关指令)
5. x到达 (周期300)
6. 执行 y = x * 2 (周期301)
结果:300周期延迟被完全隐藏!
硬件成本:
| 组件 | 作用 | 晶体管成本 |
|---|---|---|
| 重排序缓冲区(ROB) | 存储"执行中但未提交"的指令 | ~1000 万 |
| - 容量 | 200-300 条指令 | |
| - 每条存储 | 源/目标操作数、状态、异常信息 | |
| 寄存器重命名 | 消除"假依赖" | ~500 万 |
| - 架构寄存器 | 16-32 个(ISA 定义) | |
| - 物理寄存器 | 200+个(实际硬件) | |
| 依赖检查逻辑 | 判断指令间依赖 | ~500 万 |
总计:约 2000 万晶体管,占单核的 25%。
3.3 技术 2:分支预测 (Branch Prediction)
程序充满了if-else语句。CPU 必须在条件判断完成前就猜测执行哪条路径。
if (x > 0) { // 条件判断需要时间
y = x * 2; // 路径A
} else {
y = x + 10; // 路径B
}
z = y + 5; // 汇合点
分支预测器的演进:
| 代际 | 技术 | 准确率 | 失败代价 |
|---|---|---|---|
| 早期 | 静态预测(总是跳转) | 60-70% | 15-20 周期 |
| Pentium | 局部历史表 | 85-90% | 15-20 周期 |
| 现代 | 全局+局部+感知器 | 95-98% | 15-20 周期 |
直白版:2%失败的真实代价
假设一个典型程序(100 万条指令):
程序分析:
├─ 总指令数:100万条
├─ 其中分支指令:20万条(20%是if/for/while等)
└─ 分支预测准确率:98%
计算失败次数:
20万 × 2% = 4000次预测失败
计算性能损失:
• 预测成功的开销:20万 × 1周期 = 20万周期
• 预测失败的惩罚:4000 × 20周期 = 8万周期
• 总开销:20万 + 8万 = 28万周期
如果预测100%准确:
• 总开销:20万周期
实际性能损失:
8万 / 28万 = 28.6%
关键洞察:即使 98%准确率,剩下 2%的失败仍然吃掉近 30%的性能!这就是为什么现代 CPU 愿意用 2000 万晶体管做分支预测器。
硬件成本:
- 全局历史寄存器
- 局部历史表(每个分支指令)
- 混合预测器(机器学习算法)
- 总计:~2000 万晶体管
3.4 技术 3:多级缓存 (Cache Hierarchy)
CPU 通过缓存减少访问主存的次数:
访问延迟 (Intel Core i9-13900K @ 2.5GHz):
L1 Cache (32KB/core) ████ 4周期 (1.6ns)
L2 Cache (1MB/core) ████████████ 12周期 (4.8ns)
L3 Cache (36MB shared) ████████████████████████████████████████ 40周期 (16ns)
Main Memory (DDR5) ███████████████████████████████████████████████████████████████████████ 300周期 (120ns)
未命中的代价:
L1 miss → L2: 3×延迟
L2 miss → L3: 10×延迟
L3 miss → DRAM: 75×延迟
缓存的晶体管成本:
| 缓存级别 | 容量 | 晶体管数 | 占比 |
|---|---|---|---|
| L1 (I+D) | 32KB×2 | ~5000 万 | 6% |
| L2 | 1MB | ~1 亿 | 12% |
| L3 | 36MB | ~30 亿 | 22% |
| 总计 | - | ~32 亿 | 40% |
So What:40%的晶体管用于缓存,只为了减少 300 周期的主存延迟。
3.5 CPU 晶体管预算总结
现在我们可以完整理解 CPU 的 90%非计算晶体管去哪了:
CPU单核晶体管分配 (总计8亿)
目标:最小化单线程延迟
│
├─ 缓存 (40%, 3.2亿)
│ └─ 减少内存访问次数
│ └─ L3 miss代价:300周期
│
├─ 乱序执行 (25%, 2亿)
│ └─ 隐藏指令依赖
│ └─ 可同时跟踪200+条指令
│
├─ 分支预测 (10%, 0.8亿)
│ └─ 避免流水线停顿
│ └─ 失败代价:15-20周期
│
├─ 其他控制 (15%, 1.2亿)
│ └─ 译码、调度、总线等
│
└─ 计算单元 (10%, 0.8亿) ← 实际做运算的部分
└─ ALU、FPU、向量单元
关键洞察:CPU 是"延迟优化的极致"——为了让单个线程跑得快,不惜牺牲 90%的硬件资源。
过渡到 GPU:现在我们理解了 CPU 的做法。GPU 选择了完全相反的路径——不试图隐藏单个线程的延迟,而是运行如此多的线程,使得总有线程在执行。让我们看看 GPU 如何用 5%的控制逻辑管理数千个线程。
第四章:GPU 的吞吐量最大化——简单但高效
4.1 核心策略:用线程数量隐藏延迟
GPU 面对相同的 300 周期内存延迟,但采用了完全不同的策略:不尝试隐藏单个线程的延迟,而是运行足够多的线程,确保调度器总能找到就绪的线程执行。
延迟隐藏的数学原理
假设:
- 内存延迟:400 周期
- 执行一个 warp(32 线程)的一条指令:4 周期
需要多少个 warp 才能隐藏延迟?
需要的warp数 = ceil(内存延迟 / warp执行时间)
= ceil(400 / 4)
= 100 warp
每个warp有32线程:
100 warp × 32 threads = 3200个活跃线程
关键优势:零开销的上下文切换
- CPU 切换线程:需要保存/恢复寄存器、刷新 TLB (数百周期)
- GPU 切换 warp:所有 warp 的上下文同时存在于硬件中,切换只需改变一个指针 (0 周期)
直白版:工厂机器的类比
想象一台 CNC 机床(对应一个 CUDA Core)加工零件:
CPU 的策略(单订单优化):
订单1的流程:
├─ 装夹工件:10秒
├─ 等待工具到位:60秒(机器空转!)
├─ 加工:10秒
└─ 卸下工件:10秒
总时间:90秒
机器利用率:30秒工作 / 90秒 = 33%
GPU 的策略(多订单并行):
同时接10个订单,轮流使用机器:
时间轴:
0-10秒: 订单1装夹
10-20秒: 订单2装夹(订单1在等工具)
20-30秒: 订单3装夹(订单1,2都在等工具)
...
90-100秒: 订单10装夹
100-110秒:订单1工具到位了!开始加工
110-120秒:订单2工具到位了!开始加工
...
机器利用率:接近100%(永远有订单在执行某个步骤)
关键洞察:
- CPU:1 个订单 × 90 秒 = 利用率 33%
- GPU:10 个订单轮流 → 利用率 100%
- 等待时间(60 秒)被其他订单的工作填满了
映射到 GPU:
- 机器 = CUDA Core
- 订单 = warp
- 等工具 = 等内存数据
- 10 个订单 = 100 个 warp(实际 H100 有 64 个 warp/SM)
实际硬件配置 (H100 单个 SM)
| 资源 | 容量 | 支持的并发度 |
|---|---|---|
| CUDA Core | 128 个 | - |
| 寄存器堆 | 256KB (65,536 个寄存器) | 每线程 32-64 个寄存器 |
| 最大 warp 数 | 64 | 64×32 = 2048 线程 |
| 最大线程数 | 2048 | - |
| Shared Memory | 228KB | 线程块间共享 |
计算示例:
- 如果每线程使用 32 个寄存器
- 可同时支持:65,536 / 32 = 2048 个线程 ✓
- 组织为:2048 / 32 = 64 个 warp
So What:2048 个线程同时驻留在硬件中,足以隐藏任何内存延迟。
4.2 Streaming Multiprocessor (SM):GPU 的基本单元
GPU 不是简单地"塞了一堆 CUDA Core"。它需要一种方式组织这些核心、管理执行、提供内存。这就是 SM 的作用。
SM 的设计动机
| 动机 | 解决方案 | 效果 |
|---|---|---|
| 资源共享 | 多个 CUDA Core 共享调度器、内存 | 减少复制成本 |
| 可扩展性 | 通过复制 SM 线性扩展性能 | 从 108 个 SM 到 132 个 SM |
| 隔离性 | 不同 SM 独立运行 | 一个 SM 停顿不影响其他 |
类比:SM 像一个小型"GPU 芯片",完整 GPU 是 132 个"小芯片"通过 L2 Cache 连接。
H100 SM 的物理布局
SM结构 (面积约5mm², 6亿晶体管)
┌─────────────────────────────────────────────┐
│ CUDA Core阵列 (60%面积, 3.6亿晶体管) │
│ │
│ Sub-Partition 0 Sub-Partition 1 ... │
│ ┌─────────────┐ ┌─────────────┐ │
│ │ 32×FP32 Core│ │ 32×FP32 Core│ │
│ │ 16×INT32 │ │ 16×INT32 │ │
│ └─────────────┘ └─────────────┘ │
│ │
│ 共4个Sub-Partition = 128 CUDA Core │
├─────────────────────────────────────────────┤
│ 寄存器堆 (10%面积, 0.6亿晶体管) │
│ ┌─────────────────────────────────────────┐│
│ │ 256KB = 65,536个32-bit寄存器 ││
│ │ 划分为32个Bank,支持高并发访问 ││
│ └─────────────────────────────────────────┘│
├─────────────────────────────────────────────┤
│ Shared Memory / L1 Cache (8%, 0.48亿) │
│ 228KB SRAM,可动态配置比例 │
├─────────────────────────────────────────────┤
│ Warp调度器 × 4 (5%面积, 0.3亿晶体管) │
│ ┌────┐ ┌────┐ ┌────┐ ┌────┐ │
│ │Sch0│ │Sch1│ │Sch2│ │Sch3│ │
│ └────┘ └────┘ └────┘ └────┘ │
├─────────────────────────────────────────────┤
│ Tensor Core × 4 (15%, 0.9亿) │
│ Load/Store Units × 32 │
│ Special Function Units × 16 │
└─────────────────────────────────────────────┘
关键观察:
- 60%面积用于计算 (vs CPU 的 10%)
- 调度器只占 5% (vs CPU 的 35%)
- 寄存器堆很大 (256KB vs CPU 的几 KB)
4.3 CUDA Core:简单的计算单元
重要澄清:CUDA Core ≠ CPU Core
| 特性 | CPU Core | CUDA Core |
|---|---|---|
| 定义 | 完整的处理器 | 单个 ALU 单元 |
| 独立性 | 有自己的 PC、译码器 | 共享 SM 的控制逻辑 |
| 复杂度 | 8 亿晶体管 | ~3 万晶体管 |
| 功能 | 可独立运行程序 | 只执行算术运算 |
CUDA Core 的内部结构
CUDA Core (Hopper架构)
输入 (从寄存器堆)
│
├─ FP32浮点通路
│ │
│ ├─ Step 1: 乘法 (a × b)
│ │ └─ 尾数乘法器 (24-bit × 24-bit)
│ │ └─ Wallace Tree结构
│ │
│ ├─ Step 2: 加法 ((a×b) + c)
│ │ ├─ 对齐指数
│ │ ├─ 尾数加法 (48-bit)
│ │ ├─ 规范化
│ │ └─ 舍入 (IEEE 754)
│ │
│ └─ 关键:单次舍入 (FMA优势)
│
└─ INT32整数通路 (并行)
└─ 32-bit ALU
└─ 加减、位运算、移位、比较
输出 (写回寄存器堆)
设计权衡:
| 方面 | CPU Core | CUDA Core | 理由 |
|---|---|---|---|
| 流水线深度 | 15-20 级 | 4 级 | GPU 不在乎单线程延迟 |
| 频率 | 5 GHz | 1.8 GHz | 浅流水线限制频率 |
| 分支失败代价 | 15-20 周期 | 4 周期 | 浅流水线优势 |
| FP32/INT32 | 共享单元 | 独立通路 | GPU 提升吞吐量 |
关键创新:Turing 架构(2018)引入独立 INT32 通路
// 典型的CUDA代码
int idx = blockIdx.x * blockDim.x + threadIdx.x; // INT32计算地址
float value = data[idx]; // FP32读取
float result = value * 2.0f; // FP32计算
如果 INT32 和 FP32 共享单元,一半时间 CUDA Core 会闲置。独立通路使吞吐量翻倍。
4.4 Warp 调度器:简单但高效
CPU 的乱序执行引擎需要 2 亿晶体管。GPU 的 warp 调度器只需50 万晶体管(1/400),却能高效管理 64 个 warp。
为什么 GPU 调度器可以如此简单?
原因 1:编译时调度
NVIDIA 的 PTX 编译器会重排指令,将有依赖的指令分开:
原始代码:
r1 = a + b // 周期0
r2 = r1 * 2 // 依赖r1
编译器优化后:
r1 = a + b // 周期0
r3 = c + d // 周期0 (无依赖,填充空隙)
r4 = e * f // 周期1
r2 = r1 * 2 // 周期4 (r1已就绪)
CPU 需要硬件动态寻找无依赖指令,GPU 在编译时就完成了。
原因 2:warp 级并行
即使单个 warp 停顿,调度器立即切换到另一个 warp。不需要在单个线程内寻找指令级并行。
Warp 调度算法
每个时钟周期,每个调度器:
1. 遍历所有warp (64个)
└─ 跳过未就绪的warp
├─ 等待内存?
├─ 等待屏障(__syncthreads)?
└─ 等待依赖?
2. 检查资源冲突
└─ 需要的执行单元是否空闲?
3. 选择一个warp
└─ 策略:Greedy-Then-Oldest (GTO)
└─ 优先调度最老的warp
└─ 更快释放资源(寄存器、Shared Mem)
4. 发射指令到执行单元
硬件成本对比:
| 组件 | CPU 乱序引擎 | GPU warp 调度器 | 比率 |
|---|---|---|---|
| 晶体管数 | 2 亿 | 50 万 | 400× |
| 功能 | 动态重排、推测执行 | 简单选择就绪 warp | - |
| 跟踪指令数 | 200-300 条 | 64 个 warp PC | - |
| 延迟隐藏 | 硬件推测 | 线程切换 | - |
4.5 寄存器堆:海量线程的状态存储
GPU 需要同时存储 2048 个线程的寄存器状态,这是 CPU 的 1000 倍。
规模对比
| 维度 | CPU (单核) | GPU (单 SM) | 比率 |
|---|---|---|---|
| 并发线程数 | 2 (SMT) | 2048 | 1024× |
| 架构寄存器/线程 | 16 (x86-64) | 255 | 16× |
| 物理寄存器池 | 200 | 65,536 | 327× |
| 寄存器堆大小 | 1.6 KB | 256 KB | 160× |
访问带宽挑战
每个时钟周期,4 个 warp 调度器可能同时发射 4 条 FMA 指令:
带宽需求:
读操作:4 warp × 32 thread × 3 operand = 384次读/周期
写操作:4 warp × 32 thread × 1 result = 128次写/周期
总计:512次访问/周期
如何用 SRAM(每单元只有 1-2 个端口)实现?
解决方案:多体交叉存储 (Multi-Banking)
将65,536个寄存器划分为32个Bank
┌──────┬──────┬──────┬─────┬──────┬──────┐
│Bank 0│Bank 1│Bank 2│ ... │Bank30│Bank31│
│ 2KB │ 2KB │ 2KB │ │ 2KB │ 2KB │
└──────┴──────┴──────┴─────┴──────┴──────┘
地址映射:
物理地址 = (寄存器编号 × 32 + 线程ID) % 65536
为什么有效?
• Warp的32个线程访问相同寄存器号(如R5)
• 但映射到32个不同的Bank
• 可以并行访问
示例:
- 线程 0 访问 R5 → Bank (5×32+0) % 32 = Bank 0
- 线程 1 访问 R5 → Bank (5×32+1) % 32 = Bank 1
- ...
- 线程 31 访问 R5 → Bank (5×32+31) % 32 = Bank 31
代价:
- 复杂的地址生成逻辑
- 32×32 交叉开关(Crossbar)
- 约占寄存器堆总晶体管的 20%
4.6 GPU 晶体管预算总结
现在我们理解了 GPU 如何用 5%的控制逻辑管理 65%的计算单元:
GPU SM晶体管分配 (总计6亿)
目标:最大化吞吐量
│
├─ CUDA Core (60%, 3.6亿)
│ └─ 128个简单计算单元
│ └─ FP32 + INT32并行
│
├─ Tensor Core (15%, 0.9亿)
│ └─ 专用矩阵运算
│
├─ 寄存器堆 (10%, 0.6亿)
│ └─ 256KB,存储2048线程状态
│ └─ 32-way Banking
│
├─ Shared Memory/L1 (8%, 0.48亿)
│ └─ 228KB,显式管理
│
├─ 调度器 (5%, 0.3亿) ← 简单的控制逻辑
│ └─ 4个warp调度器
│ └─ 简单的就绪检查+选择
│
└─ 其他 (2%, 0.12亿)
关键洞察:GPU 通过"简化控制+海量并行"实现了比 CPU 高 6 倍的计算密度。
第五章:为什么深度学习完美匹配 GPU?
理解了 GPU 的架构后,现在可以回答:为什么深度学习能如此完美地发挥 GPU 的优势?
5.1 神经网络的计算特征
特征 1:高 Arithmetic Intensity (AI)
让我们分析一个典型的全连接层:
output = input @ weight + bias
维度:
input: (512, 1024) # batch_size × in_features
weight: (1024, 4096) # in_features × out_features
output: (512, 4096) # batch_size × out_features
计算量:
矩阵乘法:512 × 1024 × 4096 × 2 (乘+加) = 4.3 GFLOP
加bias: 512 × 4096 = 2.1 MFLOP (可忽略)
总计:约4.3 GFLOP
内存访问量:
读input: 512 × 1024 × 4B = 2.1 MB
读weight: 1024 × 4096 × 4B = 16.8 MB
读bias: 4096 × 4B = 16.4 KB
写output: 512 × 4096 × 4B = 8.4 MB
总计:27.3 MB
Arithmetic Intensity:
AI = FLOPs / Bytes
= 4.3 GFLOP / 27.3 MB
= 157 FLOP/Byte
直白版:厨师做菜的类比
低 AI 任务(如 Embedding 查找):
流程:
1. 去冰箱拿食材:10次(每次30秒)
2. 切菜炒菜:5分钟
3. 完成
分析:
├─ 取食材时间:10 × 30秒 = 5分钟
├─ 做菜时间:5分钟
└─ AI = 5分钟做菜 / 10次取食材 = 0.5分钟/次
瓶颈:一直在跑冰箱和厨房之间!
即使炒菜技术再好(计算快),大部分时间都浪费在"取数据"上
高 AI 任务(如矩阵乘法):
流程:
1. 去冰箱拿食材:1次(30秒)
2. 用这些食材做20道菜:100分钟
3. 完成
分析:
├─ 取食材时间:30秒
├─ 做菜时间:100分钟
└─ AI = 100分钟做菜 / 1次取食材 = 100分钟/次
瓶颈:做菜速度(计算能力)
只取一次食材,但用这些食材做大量菜品
映射到计算:
| 厨师类比 | 计算术语 | 说明 |
|---|---|---|
| 去冰箱 | 访问内存 | 延迟高(100ns) |
| 食材 | 数据 | weight 矩阵 |
| 炒菜 | 计算 | 矩阵乘法 |
| 炒菜技术 | 计算峰值 | FLOPS |
| 跑步速度 | 内存带宽 | GB/s |
为什么高 AI 对 GPU 有利?
| 处理器 | 计算峰值 | 内存带宽 | 适合场景 |
|---|---|---|---|
| CPU | 1 TFLOPS | 100 GB/s | 低 AI 任务(去冰箱快,但炒菜慢) |
| GPU | 60 TFLOPS | 3000 GB/s | 高 AI 任务(去冰箱也快,炒菜超快) |
对于低AI任务(AI < 10):
• 瓶颈在内存带宽
• GPU优势:3000/100 = 30×
对于高AI任务(AI > 100,如矩阵乘法):
• 瓶颈在计算能力
• GPU优势:60/1 = 60×
• 而且数据重用率高,内存不是瓶颈
关键洞察:神经网络的矩阵运算 AI 很高(>100),这意味着:
- 每个数据被重复使用很多次
- 计算时间远超内存访问时间
- GPU 的高计算峰值优势充分发挥
- 即使 GPU 内存带宽只有 CPU 的 30 倍,性能可以快 60 倍
为什么 AI=157 很高?
让我们用 Roofline 模型分析:
| 处理器 | 计算峰值 | 内存带宽 | 平衡点 | 对 AI=157 的 workload |
|---|---|---|---|---|
| CPU | 1 TFLOPS | 100 GB/s | 10 FLOP/Byte | 计算受限 → 接近 1TFLOPS |
| GPU | 60 TFLOPS | 3000 GB/s | 20 FLOP/Byte | 计算受限 → 接近 60TFLOPS |
关键洞察:
- 矩阵足够大时,每个数据元素被重用多次
- 计算量远超内存访问量
- GPU 虽然内存带宽只有 CPU 的 30×,但计算峰值有 60×
- 因为是计算受限,GPU 的性能优势充分发挥
特征 2:完美的数据并行性
# 训练一个batch
for sample in batch:
output = model(sample)
loss = criterion(output, target)
loss.backward()
关键观察:每个 sample 的前向/反向传播完全独立,没有跨 sample 的数据依赖。
这意味着:
| 执行方式 | 时间复杂度 | 理论加速比 |
|---|---|---|
| 串行 (CPU) | 512 × T | 1× |
| 并行 (GPU) | T (如果有足够资源) | 512× |
GPU 的 16,896 个 CUDA Core 天然适合这种 workload。
5.2 实际案例:ResNet-50 推理
让我们用实际数据验证 GPU 的优势。
模型信息
| 属性 | 值 |
|---|---|
| 参数量 | 25.6M |
| 单张图片 FLOPs | 4.1 GFLOP |
| Batch size | 64 |
| 总 FLOPs | 262 GFLOP |
性能对比
CPU (Intel Core i9-13900K, 16 核)
峰值性能:1 TFLOPS
实际性能:~400 GFLOPS (考虑缓存效率)
Batch 64推理时间:
262 GFLOP / 400 GFLOPS = 655 ms
吞吐量:64 / 0.655 = 98 images/s
GPU (NVIDIA H100)
峰值性能:60 TFLOPS (FP32)
实际性能:~40 TFLOPS (考虑内存开销)
Batch 64推理时间:
262 GFLOP / 40000 GFLOPS = 6.6 ms
吞吐量:64 / 0.0066 = 9,697 images/s
结果总结
| 指标 | CPU | GPU | GPU 加速比 |
|---|---|---|---|
| 推理时间 | 655 ms | 6.6 ms | 99× |
| 吞吐量 | 98 img/s | 9,697 img/s | 99× |
| 能效 | 0.33 img/s/W | 32.3 img/s/W | 98× |
为什么加速比如此巨大?
GPU优势分解:
│
├─ 计算峰值:60× (60 TFLOPS vs 1 TFLOPS)
├─ 内存带宽:30× (3000 GB/s vs 100 GB/s)
├─ 并行效率:卷积操作的并行度远超CPU的16核
└─ 实际加速:~100× (三者综合)
5.3 GPU 的局限性
GPU 并非万能。以下场景 GPU 表现不佳:
场景 1:小批量推理
Batch size = 1 (单张图片)
CPU:4.1 GFLOP / 400 GFLOPS = 10.25 ms
GPU:同样需要初始化、数据传输等开销 ≈ 5 ms
只有2×加速,不值得使用GPU的复杂性
原因:GPU 需要海量并行才能隐藏延迟。单张图片并行度不够。
场景 2:复杂控制流
def dynamic_forward(x):
if x.sum() > 0:
return branch_A(x)
else:
return branch_B(x)
Warp 分支分歧 (Divergence):
Warp中32个线程:
├─ 16个线程:条件为True → 执行branch_A
└─ 16个线程:条件为False → 执行branch_B
GPU必须串行执行:
1. 执行branch_A (只有16线程活跃,50%效率)
2. 执行branch_B (另16线程活跃,50%效率)
有效性能 = 峰值 × 50% = 30 TFLOPS
CPU无此问题,每核心独立执行
场景 3:内存密集型操作
# 大型Embedding查找 (推荐系统)
embedding = EmbeddingTable[sparse_indices]
特点:
• FLOPs很少 (只是查表)
• 内存访问量大 (随机访问)
• AI < 1 FLOP/Byte
性能对比:
| 指标 | CPU | GPU | GPU 优势 |
|---|---|---|---|
| 理论带宽 | 100 GB/s | 3000 GB/s | 30× |
| 随机访问利用率 | 80% | 20% | - |
| 有效带宽 | 80 GB/s | 600 GB/s | 7.5× |
| 考虑 PCIe/Launch 开销 | - | - | 2-3× |
结论:对于低 AI 的 workload,GPU 优势大幅缩小。
5.4 适用场景总结
| 场景 | CPU 优势 | GPU 优势 | 推荐 |
|---|---|---|---|
| 大批量训练/推理 | - | ✓✓✓ | GPU |
| 小批量推理 (batch≤4) | ✓ | ✓ | CPU |
| 复杂控制流 | ✓✓ | - | CPU |
| 高 AI 操作 (矩阵乘法) | - | ✓✓✓ | GPU |
| 低 AI 操作 (Embedding) | ✓✓ | ✓ | CPU |
| 实时交互 (低延迟) | ✓✓ | - | CPU |
| 离线批处理 (高吞吐) | - | ✓✓✓ | GPU |
第六章:总结与展望
6.1 核心权衡的本质
CPU 和 GPU 代表了处理器设计空间的两个极端:
| 设计维度 | CPU (延迟优先) | GPU (吞吐量优先) |
|---|---|---|
| 单核复杂度 | 高 (乱序执行、分支预测) | 低 (顺序执行、简单调度) |
| 核心数量 | 少 (8-64 核) | 多 (10,000+核) |
| 控制逻辑 | 多 (50%晶体管) | 少 (5%晶体管) |
| 计算单元 | 少 (10%晶体管) | 多 (65%晶体管) |
| 缓存策略 | 大容量自动管理 (L3 达 36MB) | 小容量显式管理 (程序员控制) |
| 线程切换 | 昂贵 (数百周期) | 零开销 (硬件多线程) |
| 适合场景 | 串行代码、复杂控制流 | 数据并行、简单重复计算 |
| 编程模型 | 简单 (隐式并行) | 复杂 (显式并行管理) |
| 目标应用 | 操作系统、浏览器、游戏引擎 | 科学计算、深度学习、渲染 |
关键洞察:这些权衡不是任意的,而是由以下因素共同决定:
- 物理约束:功耗墙、内存墙
- 应用特征:串行 vs 并行、复杂控制 vs 简单计算
- 资源有限:晶体管/面积/功耗预算
- 设计目标:延迟 vs 吞吐量
6.2 数字总结:关键对比
| 维度 | CPU | GPU | 差异 |
|---|---|---|---|
| 核心数 | 8-64 | 10,000+ | 100-1000× |
| 单核频率 | 5 GHz | 1.8 GHz | 0.36× |
| 单线程延迟 | 优化极致 | 慢 22× | - |
| 多线程吞吐 | 基准 | 快 100-1000× | - |
| 晶体管 → 计算 | 10% | 65% | 6.5× |
| 晶体管 → 控制 | 50% | 5% | 0.1× |
| 峰值 FLOPS | 1 TFLOPS | 60 TFLOPS | 60× |
| 内存带宽 | 100 GB/s | 3000 GB/s | 30× |
| 能效 | 3 GFLOPS/W | 200 GFLOPS/W | 67× |
6.3 未来趋势:相互学习
两种架构正在相互借鉴:
CPU 学习 GPU
| 技术 | 示例 | 效果 |
|---|---|---|
| 矢量指令 | AVX-512 | 一次处理 16 个 FP32 |
| 更多核心 | AMD Threadripper 64 核 | 提升并行能力 |
| 片上加速器 | Intel AMX (矩阵扩展) | 专用 AI 加速 |
限制:仍需保持单线程性能,控制逻辑占比难以大幅降低。
GPU 学习 CPU
| 技术 | 示例 | 效果 |
|---|---|---|
| 更大缓存 | H100 的 50MB L2 | 减少 DRAM 访问 |
| 复杂控制流 | 动态并行、条件执行 | 支持更多算法 |
| 统一内存 | CUDA Unified Memory | 简化编程 |
限制:仍需保持海量并行优势,控制逻辑不能过度复杂。
混合方案
未来处理器形态:
CPU-GPU融合
│
├─ Apple统一内存架构
│ └─ CPU和GPU共享内存,零拷贝
│
├─ NVIDIA Grace Hopper
│ └─ CPU芯片+GPU芯片,900GB/s NVLink互联
│
└─ 异构多核
└─ 大核(CPU风格) + 小核(GPU风格) + 专用加速器
└─ 根据workload动态分配
6.4 五个关键要点
-
功耗墙终结了频率竞赛
- P ∝ f³ 使得 5GHz 以上不可行
- 迫使处理器走向多核和专用化
-
CPU 牺牲 90%晶体管用于控制
- 目标:最小化单线程延迟
- 代价:计算密度低
-
GPU 牺牲单线程性能换取海量并行
- 简化控制逻辑(5%) → 更多计算单元(65%)
- 用线程数量隐藏延迟
-
深度学习完美匹配 GPU
- 高 Arithmetic Intensity (AI>100)
- 完美数据并行性 (batch>100)
- 简单控制流
-
没有绝对的"更好",只有"更适合"
- CPU:串行代码、复杂逻辑、低延迟
- GPU:并行计算、简单重复、高吞吐
最后的思考: CPU 和 GPU 的分野不是技术优劣,而是不同问题需要不同工具。 就像锤子和螺丝刀都是好工具,但你不会用锤子拧螺丝。理解硬件设计的权衡,才能在正确的场景选择正确的架构。