GPU 与并行计算

2012 年，多伦多大学的 AlexNet 在 ImageNet 图像识别大赛中，将前五误差率从 26% 降到 15.3%——领先第二名超过 10 个百分点。震撼行业的不只是准确率，还有它的实现方式：用两块 NVIDIA GTX 580 GPU（每块 3GB 显存）训练出来。

这是深度学习革命的奇点时刻。GPU，这个原本为游戏渲染而生的芯片，成为人工智能时代最重要的计算基础设施。NVIDIA 的市值从 2012 年的约 70 亿美元，飙升至 2024 年超过 3 万亿美元。

破除误解：GPU 不是"更强的 CPU"#

许多人把 GPU 理解为"更多核心的 CPU"。这是对两种架构根本差异的误解。

CPU 和 GPU 是面向完全不同问题的两种解决方案：

CPU（Central Processing Unit）：少量强大的核心（4-128 个），每个核心有深度流水线、大缓存、复杂分支预测，擅长串行的、分支密集的复杂任务。目标是最小化单条指令的延迟（Latency）。

GPU（Graphics Processing Unit）：大量简单的核心（数千到数万个），每个核心功能有限，缓存小，无复杂分支预测，擅长对大量数据执行相同操作的任务。目标是最大化数据的吞吐量（Throughput），以并行隐藏延迟。

简单比喻：CPU 是少数几位能解任何题的天才教授；GPU 是数以千计只会做乘法的小学生，但同时做乘法。

并行的分类法：Flynn、SIMD 与 SIMT#

要说清 GPU 的"并行"到底是哪一种并行，先看一张经典地图。1966 年，Michael Flynn 按"指令流"和"数据流"的数量，把计算机分成四类（Flynn 分类法，1972 年又作扩充）：

• SISD（单指令单数据）：传统冯·诺依曼单核 CPU，一条指令处理一份数据。
• SIMD（单指令多数据）：一条指令同时作用于一批数据，如 CPU 的 SSE / AVX 向量指令。
• MISD（多指令单数据）：几乎没有真实硬件实现，多作为理论占位。
• MIMD（多指令多数据）：多核 CPU、计算集群，每个核心独立执行不同指令、处理不同数据。

GPU 既不完全是 SIMD，也不完全是 MIMD。NVIDIA 为它造了一个词：SIMT（Single Instruction, Multiple Threads）。两者的区别很关键：

• SIMD 是"一个线程发一条向量指令，作用在一条定长数据向量上"——向量宽度要由程序员显式管理。
• SIMT 是"许多独立线程各自执行同一条指令"——每个线程有自己的程序计数器和寄存器，可以读写不同地址、走不同的分支路径。

换句话说，SIMD 把并行暴露给程序员（你得手写向量化），SIMT 把并行藏在线程抽象背后（你只写单个线程的代码，硬件负责把成千上万个线程批量调度）。这种抽象让 GPU 编程门槛大降，但代价藏在后面：当同一批线程真的走了不同分支，SIMT 的"独立"只是假象，硬件仍要逐条路径串行跑完（见下文 Warp 与分支发散）。

GPU 的起源：图形渲染的数学#

GPU 的历史根植于实时三维图形。渲染一帧游戏画面，需要对数百万个三角形的每个顶点做相同的矩阵变换，对数百万个像素做相同的着色计算。这种"同样的操作作用于大量数据"的模式，完美匹配了大规模并行架构。

1999 年，NVIDIA 发布 GeForce 256，首次使用"GPU"这个词。早期 GPU 是固定功能管线；2001 年的 GeForce 3 引入了可编程顶点着色器，开始允许程序员自定义 GPU 计算。

2006 年底，NVIDIA 随 G80 架构（GeForce 8800）发布了 CUDA（Compute Unified Device Architecture）编程框架（首个公开工具包于 2007 年发布），让开发者可以用 C 语言扩展编写通用程序在 GPU 上运行，GPU 从此从图形处理器变成了通用并行计算平台——称为 GPGPU（General-Purpose GPU Computing）。

CUDA 编程模型：SIMT#

CUDA 的核心概念是 SIMT（Single Instruction, Multiple Threads）：

• 程序员写 kernel 函数（GPU 上运行的函数）
• 每次启动 kernel 时指定线程数量（如 100 万个线程）
• 所有线程执行相同的指令，但作用于不同的数据
• 线程组织成线程块（Block），线程块组织成网格（Grid）

// GPU kernel：并行对数组每个元素乘以 2
__global__ void doubleArray(float *a, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        a[i] = a[i] * 2.0f;
    }
}

// 启动 100 万个线程 doubleArray<<<1000, 1024>>>(deviceArray, 1000000); ```

GPU 内部，32 个线程组成一个 Warp（NVIDIA 术语，AMD 叫 Wavefront），Warp 内的所有线程同步执行相同指令。如果 Warp 内部出现分支（if-else），NVIDIA 使用谓词执行：两条路径都执行，不相关线程的结果被屏蔽。这使得分支在 GPU 上代价极高——"避免分支"是 GPU 编程的重要优化原则。

CUDA 不是唯一选择：OpenCL 与开放标准#

CUDA 只能跑在 NVIDIA 自家硬件上。为打破这种绑定，由 Apple 发起、提交给行业联盟 Khronos Group 的 OpenCL（Open Computing Language） 应运而生：1.0 规范于 2008 年 12 月获批，2009 年随 Mac OS X Snow Leopard 首发。OpenCL 的目标是"一次编写、跨设备运行"——同一份代码可以跑在 NVIDIA、AMD、Intel 的 GPU，甚至 CPU、FPGA 上。

理想很美，权衡却很现实：开放标准要照顾各家硬件，往往落后于 CUDA 暴露最新特性（如 Tensor Core）的节奏，极致调优也更难。结果是深度学习的科研与工业界至今高度依赖 CUDA——这正是后文"GPU 垄断与依赖"的技术根源。AMD 的 ROCm / HIP、Intel 的 oneAPI、以及 OpenAI 主导的 Triton，都是近年试图在"可移植"和"高性能"之间重新找平衡的尝试。

矩阵乘法：为什么 GPU 天生适合深度学习#

深度学习的核心计算是矩阵乘法（全连接层、注意力机制）和卷积（卷积层）。

矩阵乘法 $C = A \times B$（$A$ 是 $m \times k$，$B$ 是 $k \times n$，$C$ 是 $m \times n$）：每个输出元素 $C_{ij} = \sum_{l=1}^{k} A_{il} \cdot B_{lj}$，共有 $m \times n$ 个输出元素，每个需要 $k$ 次乘加，总操作数 $O(mnk)$。

$m \times n$ 个输出元素的计算完全独立，可以完美并行。GPU 的数千个核心正好能同时处理这些独立计算。

以 GPT-3（2020）为例：其核心的 Transformer 注意力机制每次前向传播需要约 $10^{12}$（万亿）次浮点运算，在 CPU 上以每秒 $10^{12}$ 次运算的速度约需 1 秒，而高端 GPU 的峰值性能超过 $10^{14}$ 次运算每秒，快了两个数量级。

GPU 架构的演进#

NVIDIA A100（2020，Ampere 架构）： - 6912 个 CUDA 核心 - 432 个 Tensor Core（专用于矩阵乘法，每时钟周期完成 $4 \times 4$ 矩阵乘加，比普通 CUDA 核心快 8 倍） - 80 GB HBM2e 显存，带宽 2 TB/s - 峰值 FP16 性能：312 TFLOPS

Tensor Core / Matrix Core：NVIDIA 在 Volta 架构（2017）引入的专用硬件单元，专门加速混合精度矩阵运算。AMD 称之为 Matrix Core。这是深度学习训练速度暴增的关键硬件基础。

NVLink 与 NVSwitch：允许多 GPU 之间高速直连，通信带宽远超 PCIe。NVIDIA DGX A100 包含 8 块 A100 GPU，通过 NVSwitch 互联，可作为一块"超级 GPU"使用。

并行计算的挑战#

Amdahl 定律（Amdahl's Law）：程序的并行加速比受限于串行部分。设串行部分比例为 $s$，并行部分为 $1-s$，使用 $n$ 个处理器的加速比上限为：

$$\text{Speedup} \leq \frac{1}{s + \frac{1-s}{n}}$$

当 $n \to \infty$，加速比最大为 $\frac{1}{s}$。如果程序有 5% 的部分必须串行，无论多少个 GPU 并行，加速比上限只有 20。

内存带宽瓶颈：GPU 计算速度很快，但内存访问常是瓶颈。许多运算是"内存带宽受限"（Memory-Bound）而非"计算受限"（Compute-Bound）。优化 GPU 程序往往意味着减少内存访问次数、提高缓存利用率、使用高带宽内存（HBM）。

判断一个运算到底卡在哪，有个经典工具——Roofline 模型（Williams、Waterman、Patterson，UC Berkeley，2009）。它以"算术强度"（每读写 1 字节数据所做的浮点运算数，FLOPs/Byte）为横轴、可达性能为纵轴，画出两道"屋顶"：一道是内存带宽决定的斜顶，一道是峰值算力决定的平顶。算术强度低的运算（如向量加、逐元素激活函数）撞在斜顶上，是内存受限；算术强度高的运算（如大矩阵乘）才够得到平顶，是计算受限。这解释了为什么深度学习要把许多小算子"融合"（Kernel Fusion）成大算子——减少往返显存的次数、把算术强度推高，免得昂贵的 Tensor Core 饿着等数据。

通信开销：多 GPU 并行训练时，梯度同步（All-Reduce）的通信开销随 GPU 数量增加，可能成为主要瓶颈。

误区：GPU 并不能让一切都变快#

GPU 是吞吐机器，不是万能加速器。一个任务能不能从 GPU 获益，取决于它有没有大量彼此独立、做相同操作的工作。下面几类任务在 GPU 上往往不快，甚至更慢：

• 强串行依赖：每一步都要等上一步结果（如递归、状态机、指针追逐式的图遍历），没有可并行的宽度，数千个核心只能闲着。
• 分支发散严重：同一个 Warp 内线程走向天差地别的代码路径，硬件被迫逐条分支串行跑完，并行优势被吃光。
• 数据量太小：把数据从主机内存搬到显存、再启动 kernel 本身就有固定开销；任务太小，搬运和启动的时间会盖过计算节省下来的时间。

这正呼应 Amdahl 定律：决定一段程序能不能"上 GPU"的，往往不是它的计算量有多大，而是其中有多大比例真正可并行。

代价与争议#

能耗：NVIDIA H100 GPU 的 TDP（热设计功耗）超过 700 瓦。一个 AI 训练集群可能包含数万块 GPU，总功耗相当于一个中小城市。据 IEA 估计，AI 的能耗增长是全球数据中心能耗增速的主要驱动力。

GPU 垄断与依赖：CUDA 生态高度绑定 NVIDIA，切换到 AMD 或 Intel GPU 需要移植大量优化代码。Google TPU、Meta MTIA、特斯拉 Dojo 等自研 AI 芯片都是试图打破这种依赖的尝试。

AI 芯片军备竞赛：英伟达、AMD、Intel、谷歌、亚马逊、苹果、各国政府都在投入巨资开发新一代 AI 加速芯片，这场竞赛的走向将决定 AI 计算的成本和可及性。

跨域连接#

• 线性代数：决定 GPU 快慢的不是运算总量，而是算术强度——每读写一字节数据能做多少次浮点运算。稠密矩阵乘的运算量按三个维度相乘增长，数据量只按两两之和增长，强度随规模上升，才够得着 Roofline 的算力平顶；逐元素激活与向量加的强度是常数，永远撞在带宽斜顶上。算子融合有效的原因正在于此。
• CPU 流水线：两类架构对同一问题给出相反答案。CPU 靠乱序执行与分支预测缩短单条指令的延迟；GPU 不缩短延迟，而是准备远超核心数的线程，在一批线程等访存时切到另一批，用并行度把延迟盖住。推论可检验：一旦同一 warp 内线程走向不同分支，可切换的同质工作消失，隐藏机制随之失效。
• 半导体物理：并行不是审美偏好而是散热的后果。芯片功耗随频率与电压升高而快速增长，单位面积能散掉的热量却有上限，于是提频这条路在某点被封死，性能只能改由增加同时工作的单元来换。这既解释了为何某个年代之后是「更多核」取代「更高频」，也解释了 GPU 为何用大量简单核心。
• 网络效应：CUDA 的护城河不在指令集而在生态。库、算子、教程与已调优代码是互补品，其价值随使用者数量上升，迁移成本则随自有代码量上升。推论是开放标准即便在峰值性能上追平也难扭转局面，因为用户比较的从来不是单点性能，而是整套工具链的迁移代价。
• 碳预算与净零：单卡数百瓦、集群数万卡，训练设施的功率以数十兆瓦计，约束因此从「能买到多少芯片」转向「能拿到多少并网容量与冷却水」。两条推论随之可测：算力扩张会先撞上电网排队而非晶圆产能；选址向电力富余、气候凉爽处集中，算力的碳强度由当地电源结构决定。

参考文献#

• Kirk, D. & Hwu, W. Programming Massively Parallel Processors: A Hands-on Approach. Morgan Kaufmann, 2022.
• Nickolls, J. et al. Scalable Parallel Programming with CUDA. Queue 6(2), 2008. （CUDA 的奠基性论文）
• Amdahl, G. Validity of the Single Processor Approach to Achieving Large Scale Computing Capabilities. AFIPS, 1967.
• Flynn, M. J. Very High-Speed Computing Systems. Proceedings of the IEEE 54(12), 1966. （Flynn 分类法的提出）
• Williams, S., Waterman, A., Patterson, D. Roofline: An Insightful Visual Performance Model for Multicore Architectures. Communications of the ACM 52(4), 2009.
• Krizhevsky, A., Sutskever, I., Hinton, G. ImageNet Classification with Deep Convolutional Neural Networks. NIPS, 2012. （AlexNet 论文）
• Khronos Group. The OpenCL Specification (v1.0). 2008. （OpenCL 开放标准）

延伸阅读#

• NVIDIA. CUDA C++ Programming Guide. docs.nvidia.com/cuda/. （官方文档）