Keyboard shortcuts

Press or to navigate between chapters

Press S or / to search in the book

Press ? to show this help

Press Esc to hide this help

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 优势
CPU0.1 微秒4,167 微秒-
GPU2.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%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%
L21MB~1 亿12%
L336MB~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 Core128 个-
寄存器堆256KB (65,536 个寄存器)每线程 32-64 个寄存器
最大 warp 数6464×32 = 2048 线程
最大线程数2048-
Shared Memory228KB线程块间共享

计算示例

  • 如果每线程使用 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                 │
└─────────────────────────────────────────────┘

关键观察

  1. 60%面积用于计算 (vs CPU 的 10%)
  2. 调度器只占 5% (vs CPU 的 35%)
  3. 寄存器堆很大 (256KB vs CPU 的几 KB)

4.3 CUDA Core:简单的计算单元

重要澄清:CUDA Core ≠ CPU Core

特性CPU CoreCUDA 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 CoreCUDA Core理由
流水线深度15-20 级4 级GPU 不在乎单线程延迟
频率5 GHz1.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)20481024×
架构寄存器/线程16 (x86-64)25516×
物理寄存器池20065,536327×
寄存器堆大小1.6 KB256 KB160×

访问带宽挑战

每个时钟周期,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 有利?

处理器计算峰值内存带宽适合场景
CPU1 TFLOPS100 GB/s低 AI 任务(去冰箱快,但炒菜慢)
GPU60 TFLOPS3000 GB/s高 AI 任务(去冰箱也快,炒菜超快)
对于低AI任务(AI < 10):
• 瓶颈在内存带宽
• GPU优势:3000/100 = 30×

对于高AI任务(AI > 100,如矩阵乘法):
• 瓶颈在计算能力
• GPU优势:60/1 = 60×
• 而且数据重用率高,内存不是瓶颈

关键洞察:神经网络的矩阵运算 AI 很高(>100),这意味着:

  1. 每个数据被重复使用很多次
  2. 计算时间远超内存访问时间
  3. GPU 的高计算峰值优势充分发挥
  4. 即使 GPU 内存带宽只有 CPU 的 30 倍,性能可以快 60 倍

为什么 AI=157 很高?

让我们用 Roofline 模型分析:

处理器计算峰值内存带宽平衡点对 AI=157 的 workload
CPU1 TFLOPS100 GB/s10 FLOP/Byte计算受限 → 接近 1TFLOPS
GPU60 TFLOPS3000 GB/s20 FLOP/Byte计算受限 → 接近 60TFLOPS

关键洞察

  1. 矩阵足够大时,每个数据元素被重用多次
  2. 计算量远超内存访问量
  3. GPU 虽然内存带宽只有 CPU 的 30×,但计算峰值有 60×
  4. 因为是计算受限,GPU 的性能优势充分发挥

特征 2:完美的数据并行性

# 训练一个batch
for sample in batch:
    output = model(sample)
    loss = criterion(output, target)
    loss.backward()

关键观察:每个 sample 的前向/反向传播完全独立,没有跨 sample 的数据依赖。

这意味着:

执行方式时间复杂度理论加速比
串行 (CPU)512 × T
并行 (GPU)T (如果有足够资源)512×

GPU 的 16,896 个 CUDA Core 天然适合这种 workload。

5.2 实际案例:ResNet-50 推理

让我们用实际数据验证 GPU 的优势。

模型信息

属性
参数量25.6M
单张图片 FLOPs4.1 GFLOP
Batch size64
总 FLOPs262 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

结果总结

指标CPUGPUGPU 加速比
推理时间655 ms6.6 ms99×
吞吐量98 img/s9,697 img/s99×
能效0.33 img/s/W32.3 img/s/W98×

为什么加速比如此巨大?

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

性能对比

指标CPUGPUGPU 优势
理论带宽100 GB/s3000 GB/s30×
随机访问利用率80%20%-
有效带宽80 GB/s600 GB/s7.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)小容量显式管理 (程序员控制)
线程切换昂贵 (数百周期)零开销 (硬件多线程)
适合场景串行代码、复杂控制流数据并行、简单重复计算
编程模型简单 (隐式并行)复杂 (显式并行管理)
目标应用操作系统、浏览器、游戏引擎科学计算、深度学习、渲染

关键洞察:这些权衡不是任意的,而是由以下因素共同决定:

  1. 物理约束:功耗墙、内存墙
  2. 应用特征:串行 vs 并行、复杂控制 vs 简单计算
  3. 资源有限:晶体管/面积/功耗预算
  4. 设计目标:延迟 vs 吞吐量

6.2 数字总结:关键对比

维度CPUGPU差异
核心数8-6410,000+100-1000×
单核频率5 GHz1.8 GHz0.36×
单线程延迟优化极致慢 22×-
多线程吞吐基准快 100-1000×-
晶体管 → 计算10%65%6.5×
晶体管 → 控制50%5%0.1×
峰值 FLOPS1 TFLOPS60 TFLOPS60×
内存带宽100 GB/s3000 GB/s30×
能效3 GFLOPS/W200 GFLOPS/W67×

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 五个关键要点

  1. 功耗墙终结了频率竞赛

    • P ∝ f³ 使得 5GHz 以上不可行
    • 迫使处理器走向多核和专用化
  2. CPU 牺牲 90%晶体管用于控制

    • 目标:最小化单线程延迟
    • 代价:计算密度低
  3. GPU 牺牲单线程性能换取海量并行

    • 简化控制逻辑(5%) → 更多计算单元(65%)
    • 用线程数量隐藏延迟
  4. 深度学习完美匹配 GPU

    • 高 Arithmetic Intensity (AI>100)
    • 完美数据并行性 (batch>100)
    • 简单控制流
  5. 没有绝对的"更好",只有"更适合"

    • CPU:串行代码、复杂逻辑、低延迟
    • GPU:并行计算、简单重复、高吞吐

最后的思考: CPU 和 GPU 的分野不是技术优劣,而是不同问题需要不同工具。 就像锤子和螺丝刀都是好工具,但你不会用锤子拧螺丝。理解硬件设计的权衡,才能在正确的场景选择正确的架构。