8.1 GPU 運算與 CUDA:算力如何驅動深度學習(2007-2012)#
一、歷史背景:思想早已就位,等待的是算力#
回顧前七章會發現一個規律:反向傳播(1986)、CNN(1989)、LSTM(1997)的核心思想都在 1998 年前就緒,但神經網路的實用化卻等到 2012 年之後。缺失的不是智慧,而是每秒一億億次乘加的原始算力。本章講述 GPU(圖形處理器)如何從「畫像素的晶片」變成「深度學習的心臟」,以及這場硬體革命如何反向塑造了神經網路架構的形狀。
timeline
title GPU 與深度學習的匯流(1999-2016)
1999 : GeForce 256(首塊標榜 GPU 的消費卡)
2001 : 可程式化著色器(頂點/像素著色器)
2006 : NVIDIA G80 + CUDA 1.0(通用 GPU 運算誕生)
2007-2009 : 學界用 CUDA 做 GPGPU(含神經網路初步實驗)
2010 : Dan Cireşan(IDSIA)用 GPU 訓 Schmidhuber 的深 CNN 奪多項競賽
2012 : AlexNet 用 2×GTX 580 奪 ImageNet(cuda-convnet, Krizhevsky)
2014 : cuDNN 發布(深度學習原語的極致優化)
2016 : NVIDIA Pascal + 專用 FP16 tensor 核(P100)二、為什麼是 GPU:神經網路計算的本質#
2.1 神經網路 = 大規模稠密線性代數#
回顧各章的計算核心:
- 前向傳播:一連串矩陣乘法 \(z = Wx + b\) (3 章);
- 卷積:可展開(im2col)為矩陣乘法(5 章);
- 反向傳播: \(W^\top \delta\) 與 \(\delta a^\top\) ——還是矩陣乘法(3.2);
- 注意力: \(QK^\top\) 、 \(\text{softmax}(\cdot)V\) ——矩陣乘法(7 章)。
深度學習的 90% 以上浮點運算是矩陣乘法(GEMM),而 GEMM 的特徵是:計算量巨大( \(O(n^3)\) 乘加)、數據重用率極高(每個元素被反覆使用)、各計算單元之間相互獨立。這正是平行計算的理想形態。
2.2 CPU 與 GPU 的設計哲學差異#
| 維度 | CPU | GPU |
|---|---|---|
| 核心數 | 數個-數十個(複雜核) | 數千個(簡單核) |
| 設計目標 | 低延遲(儘快完成單線程) | 高吞吐(儘快完成海量同質任務) |
| 電晶體用途 | 分支預測、亂序執行、大 cache | 大量 ALU |
| 適合 | 邏輯分支、串行依賴 | 稠密線性代數、SIMT 模式 |
神經網路訓練恰好是「海量獨立的乘加、幾乎沒有分支」——CPU 的電晶體花在分支預測上是浪費,GPU 的數千簡單核正中要害。量化感受:2012 年 GTX 580 的單精度峰值約 1.6 TFLOPS,同代頂級 CPU 約 0.15 TFLOPS——一個數量級的差距,且 GPU 可以插多塊。
2.3 為什麼是 CUDA 而非圖形 API#
GPU 平行計算並非新事——學界在 2004 年前後曾用圖形著色語言(OpenGL/DirectX shader)「欺騙」GPU 做通用計算,如 BrookGPU(2003)。但著色器程式模型扭曲(沒有真正的隨機寫入、缺乏通用記憶體抽象),開發極痛苦。2006-2007 年 NVIDIA 的 G80 架構與 CUDA 1.0 改變一切:
- 統一著色器架構:流式處理器(SM)不再區分頂點/像素單元,成為通用可程式化核心;
- C 語言介面:科學家不需要懂圖形學即可寫平行程式;
- 真正的記憶體模型:全域記憶體、共享記憶體(片上快取)、任意讀寫。
CUDA 的歷史意義堪比馮紐曼架構之於計算機:它把平行硬體變成普通科學家可用的程式設計平台。NVIDIA 冒了巨大商業風險(華爾街當年普遍認為 GPGPU 是不務正業),而這場豪賭最終以市值萬億回報——硬體與演算法共同演化的最佳案例。
三、CUDA 程式設計模型與 AlexNet 的算力帳#
3.1 CUDA 的核心抽象#
CUDA 模型的三層結構:執行緒(thread)→ 區塊(block)→ 網格(grid)。同一區塊內的執行緒共享片上共享記憶體,可同步;區塊間由硬體調度器自動分配到各 SM。寫法是「單執行緒的視角 + 自動的千倍複製」:
// 一個最簡單的 CUDA 核:向量加法 c = a + b(每個執行緒算一個元素)
__global__ void vec_add(const float* a, const float* b, float* c, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) c[i] = a[i] + b[i];
}
// 啟動:n 個元素由 (n+255)/256 個區塊、每區塊 256 執行緒處理
// vec_add<<<(n+255)/256, 256>>>(da, db, dc, n);
效能的關鍵不在啟動多少執行緒,而在記憶體存取模式:GPU 全域記憶體延遲極高(數百週期),必須用「合併存取(coalescing,相鄰執行緒讀相鄰位址,合併為一次事務)」與共享記憶體分攤延遲。矩陣乘法的分塊(tiling)演算法正是圍繞共享記憶體設計的——深度學習框架的性能魔法,本質是這些記憶體層級技巧。
3.2 AlexNet 的算力帳#
用 AlexNet(5.2)做一次具體的算力核算:
- 前向約 7 億 FLOPs/圖;反向約 2 倍;共約 2.1 GFLOPs/樣本更新;
- 120 萬張圖 × 90 輪 ≈ 1.08 億樣本更新 ≈ ** \(2.3 \times 10^{17}\) FLOPs**;
- 2 塊 GTX 580 有效吞吐約 0.5 TFLOPS(峰值 3.2)→ 約 5-6 天。
若用 2012 年的 CPU 叢集,同樣訓練需要數月——競賽週期內根本不可能迭代。AlexNet 的勝利公式因此可以寫成:構想(1989 的 CNN)+ 資料(2009 的 ImageNet)+ 算力(2007-2012 的 CUDA)。2012 年論文的致謝明確感謝 Krizhevsky 自寫的 cuda-convnet——那兩千行 CUDA 代碼與演算法同等重要。
3.3 之後的算力軍備競賽#
flowchart LR
A["需求: 訓練越來越大的模型"] --> B["NVIDIA: 為 DL 專門優化<br>P100→V100→A100→H100<br>Tensor Core / BF16 / NVLink"]
B --> C["框架封裝: cuDNN / NCCL<br>科學家無需寫 CUDA"]
C --> D["更大模型需要更多算力"]
D --> A
A --> E["規模定律(9.2)<br>算力成為能力的代理變數"]- cuDNN(2014):把卷積、池化、歸一化等原語做到硬體極致,框架呼叫即可,科學家從此不碰 CUDA;
- Tensor Core(2017-):專門的 \(4\times4\) 矩陣乘法電路 + 半精度(FP16/BF16),吞吐再上一個數量級——硬體開始為神經網路量身定製(而非反向);
- 多卡擴展:資料並行(每卡一份模型副本、同步梯度,NCCL 實作 all-reduce)→ 模型並行、流水線並行(9.2 的大模型時代三件套)。
值得注意的正回饋環:演算法的架構選擇開始迎合硬體。ReLU 之於 sigmoid 的勝利、Transformer 之於 RNN 的勝利,都部分因為它們更「矩陣乘法友好」、更可並行。硬體不只是跑演算法的馬車,而是參與了演算法的演化選擇。
四、程式碼實作:從感知算力差距#
4.1 CPU vs GPU 的矩陣乘法對決#
import torch, time
def bench(device, n=4096, reps=20):
a = torch.randn(n, n, device=device)
b = torch.randn(n, n, device=device)
(a @ b).item() # 預熱
t0 = time.time()
for _ in range(reps):
c = a @ b
if device != "cpu": torch.cuda.synchronize()
dt = (time.time() - t0) / reps
tflops = 2 * n**3 / dt / 1e12 # 2n³ = 每次 GEMM 的浮點運算數
return dt, tflops
dev = "cuda" if torch.cuda.is_available() else ("mps" if torch.backends.mps.is_available() else "cpu")
for d in ["cpu", dev]:
if d == dev or d == "cpu":
dt, tf = bench(d)
print(f"{d:5s}: {dt*1000:7.2f} ms / 4096² GEMM = {tf:6.2f} TFLOPS")
# 典型結果:CPU ~0.1-0.5 TFLOPS,GPU ~30-300 TFLOPS —— 兩到三個數量級
# 這個差距 × 20 年硬體進步 × 多卡擴展 = 深度學習革命的物質基礎
4.2 自訂 CUDA 核(torch 的 C++/CUDA 擴展示意)#
# 用 PyTorch 自訂運算的兩個層次:
# (1)用張量操作組合(無 CUDA 知識):
def swish(x):
return x * torch.sigmoid(x) # 三個原語,框架自動並行
# (2)真正手寫 CUDA(高階用法,示意 Triton 語言——OpenAI 讓 GPU 程式設計再次大眾化):
# import triton, triton.language as tl
# @triton.jit
# def swish_kernel(x_ptr, y_ptr, n, BLOCK: tl.constexpr):
# pid = tl.program_id(0)
# offs = pid * BLOCK + tl.arange(0, BLOCK)
# mask = offs < n
# x = tl.load(x_ptr + offs, mask=mask)
# tl.store(y_ptr + offs, x * tl.sigmoid(x), mask=mask)
# Triton(2019)以 Python 語法寫 GPU 核,是「CUDA 大眾化」的當代故事
print("Triton 示意:GPU 程式設計的門檻從 C/CUDA 降到 Python")
4.3 資料並行的梯度同步(多卡訓練的原理)#
import torch, torch.nn as nn
# 資料並行的手寫原理(單機示範;實務用 DDP):
model = nn.Linear(1000, 10)
optimizer = torch.optim.SGD(model.parameters(), lr=0.1)
def data_parallel_step(batch, n_gpus=4):
"""每 GPU 獨立算子集梯度,再 all-reduce 平均 —— DDP 的本質"""
micro = len(batch) // n_gpus
grads = []
for g in range(n_gpus): # (真實中各 GPU 並行執行)
sub = batch[g*micro:(g+1)*micro]
loss = nn.functional.cross_entropy(model(sub), sub.new_zeros(len(sub)))
loss.backward() # 各自累積梯度
grads.append([p.grad.clone() for p in model.parameters()])
model.zero_grad()
for pi, p in enumerate(model.parameters()): # ← all-reduce:梯度平均
p.grad = torch.stack([g[pi] for g in grads]).mean(0)
optimizer.step()
print("資料並行 = 複製模型 + 切資料 + all-reduce 梯度")
print("把 N 張卡的吞吐變成 N 倍 —— 大模型時代的擴展基石")
五、算力、演算法與資料的共同演化#
回望歷史,算力與深度學習不是「平台與應用」的簡單關係,而是共同演化:
- 1986-2006 的黑暗等待:BP 演算法有了,但 CPU 時代一個epoch要數天,實驗迭代慢到無法積累洞見;
- 2007-2012 的解鎖:CUDA 讓訓練提速一個數量級,ImageNet 提供了「值得提速的問題」——AlexNet 證明組合威力;
- 2012-2016 的軍備競賽:架構(更深、更大)與硬體(更專、更快)互相拉動,cuDNN/Tensor Core 把「神經網路專用晶片」變成主流;
- 2017- 的規模時代:算力成為模型能力的代理變數(scaling laws,9.2),NVIDIA 從晶片公司變成 AI 時代的基礎設施壟斷者。
方法論的教訓同樣深刻:當一個好想法(BP、CNN、Transformer)被埋沒時,先問「是想法的問題還是算力的問題」。深度學習史多次證明,等待正確的硬體比等待正確的數學更划算。
六、本節重點#
- 神經網路訓練的計算本質是大規模稠密矩陣乘法(GEMM),計算密集、數據重用高、可並行——與 GPU 的 SIMT 高吞吐設計完美匹配,與 CPU 的低延遲設計失配。
- CUDA(2006-2007)以 C 語言介面 + 統一著色器架構把平行硬體大眾化,是深度學習復興的物質前提;其核心抽象是「執行緒-區塊-網格」與記憶體層級(合併存取、共享記憶體)。
- AlexNet 的算力帳: \(2.3\times10^{17}\) FLOPs 在 2 塊 GTX 580 上約一週,CPU 則需數月——算力不是加速器,而是使實驗迭代成為可能的存在條件。
- cuDNN、Tensor Core、多卡 all-reduce 把「科學家不碰 CUDA」變成現實;硬體為神經網路專門定製,演算法也反向迎合硬體(ReLU、Transformer 的勝利部分是並行友好性的勝利)。
- 算力、資料、架構三元共振的歷史規律:缺任何一角都不會有 2012;規模時代算力更成為能力的代理變數。
七、思考問題#
- 神經網路訓練「分支少、同質性高」,請再列舉兩個與 GPU 失配的計算特徵(提示:稀疏注意力、動態控制流、圖演算法),並說明它們為何讓某些模型架構吃虧。
- AlexNet 若晚三年出現(2015 年才有 K80/cuDNN),深度學習歷史會如何改寫?請論證「思想等待硬體」與「硬體等待思想」哪個是當時的瓶頸。
- CUDA 開放 C 介面但封閉生態(綁定 NVIDIA 硬體)。請從「平台壟斷」與「開源替代」(ROCm、Triton、TPU 的 XLA)角度分析 AI 算力供應鏈的競爭格局。
- 資料並行的 all-reduce 有 Ring、Tree 等演算法,通信量 \(O(\text{模型大小})\) 與卡數無關(Ring)。當模型大到單卡放不下時,資料並行失效,需要哪些替代方案?(連接 9.2 的 ZeRO、流水線、張量並行。)
- 「演算法迎合硬體」是否威脅研究的多樣性?例如稀疏模型、類比計算、神經形態晶片在「稠密 GEMM 主義」下處於劣勢。你如何評價這種生態選擇壓力?