SICI · 化神一层 · 第二十三章

并行通天

并行计算与高性能计算 · 化神一层·并行通天

你以为GPU有数千核心就是数千倍加速。可这数千核心,喂得过数据吗?带宽是天花板——算术强度才是灵魂。

带宽之壁低语:「你既信核心越多越快,与我赌一局——这五样任务,哪一样堆了核心也不加速?」
序章

雾锁并行之门

堆核心就能线性加速吗?五道任务,哪样堆了核心也白搭

门楣上刻着四个字:"并行通天。"门内没有整洁的走廊,只有一片翻涌的紫色雾——雾的尽头,矗立一面半透明的光壁,像一堵无声的悬崖。雾凝成一个人形,通体由无数条数据流构成,但没有一条流得通——它们全部撞在光壁上,碎成浪花。

紫色雾霭深处矗立一面半透明光壁,凝成人形的数据流尽数撞在壁面碎成浪花——带宽之壁
带宽之壁 带宽之壁

「来得好。你既信'核心越多越快',与我赌一局——这五样任务,哪一样堆了核心也不加速?」

诺依曼老祖 诺依曼老祖

「别急着下注。你以为并行就是多开几个核心——先把这五样点一遍看看。」

并行加速反直觉挑战器

反直觉下注(五选一)→ 加速曲线打脸 → 雾账记账 → 天花板图谱拼合

带宽之壁逼你下注:五样任务,哪一样堆了核心也不加速?逐个点击,看加速曲线。
详解 · K23-12 并行性能分析:加速比/效率/可扩展性

加速比 S(p)=T₁/Tₚ、效率 E(p)=S(p)/p;S>p 叫超线性,通常是多核缓存合起来装下了单核装不下的工作集。台上五道任务的分野全由算术强度 AI=计算量÷访存量决定:拐点左侧性能=min(峰值,带宽×AI),与核心数无关。用 S5 那台 A100(9.7 TFLOPS、1550 GB/s、拐点 6.26)手算:向量加法 1 FLOP 换 12 Byte(2读1写 float32),AI=0.083,实得 129 GFLOPS,仅峰值 1.3%;CSR SpMV 每非零元 2 FLOP 换 8 Byte,AI=0.25,实得 388 GFLOPS,4%;归约同为 0.25。而 1024² 矩阵乘法 2n³≈2.15×10⁹ FLOP、理想流量 12n²≈1.26×10⁷ Byte,AI≈171。同芯片同核心,实得差 75 倍。问可扩展性须先分清强扩展(问题固定加核,撞 Amdahl 墙)与弱扩展(每核负载固定,撞容量与通信墙)。出处:Hennessy & Patterson 第1、4章;Hwu PMPP 第4版第5章。

第一幕 · 自信→崩塌

频率之墙 · 并行之形

P∝f³ 功耗立方爆炸 · SISD/SIMD/MISD/MIMD 四象限归位

石径尽头是一台悬空的主频仪表盘,指针停在 3.5 GHz。诺依曼老祖站在盘边,带宽之壁的光从壁面渗过来,冷冷补了一句。

诺依曼老祖 诺依曼老祖

「1974年Dennard说,频率翻倍、电压跟着降,功耗不变。2006年这条路断了。你猜,继续拉频率会怎样?」

带宽之壁 带宽之壁

「拉吧。把频率推到8 GHz。看看你的功耗账单。」

频率墙推演 · Flynn四象限

预测功耗增长 → 推频率旋钮看P∝f³ → Flynn四象限拖拽归位

频率翻倍,功耗怎么涨?
详解 · K23-01 并行计算的动机:功耗墙/频率墙/ILP极限

1974 年 Dennard 等人(IEEE JSSC)给出缩放定律:尺寸缩 k 倍、电压同比降、电场不变,功率密度守恒,频率可一路涨。动态功耗 P=α·C·V²·f,若 V 随 f 同比上升就得 P∝f³。台上仪表盘正是这条立方律:以 3.5 GHz/100 W 为基准,4 GHz→100×(4/3.5)³≈149 W,6 GHz≈504 W,8 GHz≈1194 W。2004-2006 年这条路断了:栅氧薄到几个原子层出现隧穿漏电、亚阈值摆幅压不下去、掺杂涨落让阈值电压失控——登纳德缩放失效,摩尔定律却仍在供给晶体管。同期 ILP 也边际递减:流水线、超标量、乱序窗口、分支预测、寄存器重命名(Tomasulo,1967)每加一层换来越来越少的单核性能。H&P 2017 年图灵奖演讲给出曲线:单核年增速从 1986-2003 的约 52% 跌到 2011-2015 的约 12%,再到 3% 量级。两头撞墙,只剩把晶体管换成并行度。出处:Hennessy & Patterson 第1章与第4章开篇;Harris & Harris 第7章。

详解 · K23-02 Flynn分类法:SISD/SIMD/MISD/MIMD

Michael J. Flynn 1966 年在《Proceedings of the IEEE》的《Very High-Speed Computing Systems》提出(1972 年细化),只用两个二元维度:指令流单/多 × 数据流单/多。SISD 是冯·诺依曼单机,但那是逻辑视图——现代单核靠超标量物理上同时发射多条指令。SIMD 一条指令驱动一整个向量:Cray-1(1976)等向量机、CPU 的 MMX(1997)/SSE(1999)/AVX(2011)/AVX-512(2016),以及 GPU。MISD 工程上几乎不存在,教科书常举容错表决。MIMD 是多核 CPU、SMP、集群、整个 TOP500,主流写法却是 SPMD——同一份程序在 p 个处理器上跑、靠 rank 分工,MPI 即典型。SIMT 是 NVIDIA 2006-2007 随 G80/CUDA 提出的杂交体:形式像 SIMD(一次发射、32 线程锁步),语义像 MIMD(每线程独立寄存器,Volta 起还有独立 PC)。台上四象限考的就是这条边界。出处:H&P 第3、4章。

第一幕 · 探索

通天塔 · SIMT锁步

6912核心的秘密 · warp=32线程锁步执行 · SM/CUDA核心/寄存器/共享内存

墙后是一座悬空的塔,塔身由无数微小的光点构成——每个光点是一个核心。诺依曼老祖指着塔身,带宽之壁的光渗进塔基。

诺依曼老祖 诺依曼老祖

「NVIDIA A100,108个SM,6912个CUDA核心。但它不是6912个独立CPU——它用一种叫SIMT的把戏管着它们。」

带宽之壁 带宽之壁

「6912个核心。你觉得你能拿到6912倍加速?记住这个数字。」

SIMT锁步流水线动画器

展开SM → 拖32线程进warp槽 → 锁步执行 → GPU参数归位

详解 · K23-03 SIMT(单指令多线程)模型与GPU架构基础

SIMT 由 NVIDIA 2006-2007 随 G80/CUDA 推出。执行粒度是 warp:32 线程绑成一个调度单位,同拍取同一条指令、在 32 条 lane 上锁步(台上"32 人小队迈同一步")。与 SIMD 差三处:每线程有独立寄存器状态;分歧时硬件按路径串行重放、用活跃掩码关掉无关 lane;Volta(2017)起每线程有独立 PC 与调用栈。延迟隐藏是命门:A100 每 SM 驻留上限 64 warp(2048 线程),上下文全在寄存器文件里,切换零开销。但算笔账:每 SM 4 个调度器,64÷4=16 warp 归一个管;DRAM 延迟约 400 cycles,哪怕每 warp 靠 ILP 每 2 拍发一条,16 个也只填出 32 个发射机会,差一个数量级——所以必须把数据搬进 30 cycles 的共享内存,这正是 S5 分块的动机。台上填 40 线程那格同理:硬件开 ⌈40/32⌉=2 warp、64 lane,24 条被掩码,利用率 40/64=62.5%。出处:Hennessy & Patterson 第4章;Hwu PMPP 第4版第3、4章。

详解 · K23-04 GPU硬件架构:SM/CUDA核心/寄存器文件/共享内存

A100(GA100,TSMC 7nm,2020,542 亿晶体管/826 mm²)对应台上参数表:108 个 SM(完整 die 128 个,屏蔽 20 个换良率),每 SM 四分区各 16 个 FP32 核,合计 108×64=6912——这是 lane 数,不是 6912 个独立 CPU;另有 32 个 FP64 核/SM(全芯片 3456)、4 个第三代 Tensor Core、4 个 warp 调度器。峰值可手算:FP64 向量 3456×2×1.41 GHz≈9.75 TFLOPS(标称 9.7),FP32 6912×2×1.41≈19.5,TF32 Tensor Core 156(2:4 稀疏翻倍)——S5 用的 9.7 即 FP64 向量峰值。存储三件套:每 SM 256KB 寄存器文件(65536×32 bit,零开销 warp 切换的本钱);192KB 统一 L1/共享内存(可配置,单 block 最多 164KB 共享);全芯片 40MB L2。片外 5 组 HBM2e:5×1024=5120 bit ×2.43 Gbps ÷8=1555 GB/s(本章取整 1550);80GB SXM 版 2039 GB/s;SXM 整卡 TDP 400W。出处:Hennessy & Patterson 第4章;NVIDIA A100 白皮书。

第一幕 · 掌控膨胀

线程之海 · CUDA编程

线程/线程块/网格 · threadIdx+blockIdx索引 · 启动配置<<<>>>

塔顶是一片无垠的网格——数以万计的方块整齐排列,每个方块里嵌着线程,像一片被切分的海。诺依曼老祖递给你一张空白的内核卷轴。

诺依曼老祖 诺依曼老祖

「thread、block、grid。三件事。线程是最小执行单元,块是一组能共享内存的线程,网格是所有块的集合。写一个向量加法。」

带宽之壁 带宽之壁

「写吧。写完你会觉得自己无所不能。记住这种感觉。」

CUDA 内核组装 · 向量加法

拖代码块组装内核 → 索引公式 → 启动配置(block须32倍数)

详解 · K23-05 CUDA编程模型:线程/线程块/网格层次

CUDA 三层:thread 最小执行单元,block(CTA)是能用共享内存与 __syncthreads() 配合的一组线程,grid 是全部 block。配置写作 kernel<<<gridDim, blockDim>>>(args)。全局索引只有一个公式:i = threadIdx.x + blockIdx.x * blockDim.x,末尾 if (i < N) 掩掉凑整块多出的线程。硬件映射要背:block 整体指派到某一个 SM、生命周期内不迁移;一个 SM 可驻留多个 block(A100 上限 32),所以 block 间无法同步(Volta 后才有 grid.sync())。block 须是 32 的倍数、上限 1024(warp 粒度)。台上利用率就是这道算术:block=40 时 ⌈40/32⌉=2 warp、64 lane 只装 40 线程,利用率 62.5%;叠加每 SM 最多 32 block,32×40=1280 线程,占用率 1280/2048 仍是 62.5%。反之 block=256、每 SM 驻 8 块正好 2048 线程,占用率 100%。出处:Hwu PMPP 第4版第2、3章;CUDA C++ Programming Guide。

第二幕 · 暗藏隐患

层次之塔 · CUDA内存

全局/共享/寄存器/常量 · 延迟差13倍 · 数据该放哪一层

内核卷轴亮起,塔身内部显出一座四层高塔——从顶到底依次闪烁。诺依曼老祖指着塔身。

诺依曼老祖 诺依曼老祖

「寄存器最快,共享内存次之,常量内存只读缓存,全局内存最慢——差13倍。数据放哪一层,决定你的内核跑多快。」

带宽之壁 带宽之壁

「全局内存~400 cycles,共享内存~30 cycles。你以为都一样?差13倍——这13倍就是我的领地。」

内存层次之塔 · 四级归位

拖内存卡到对应层级 → 延迟对比弹出 → 数据该放哪一层

详解 · K23-06 CUDA内存层次:全局/共享/寄存器/常量内存

内存层次要同时看延迟、带宽、作用域。寄存器线程私有、编译期分配,延迟记 0;共享内存按 block 共享,Ampere 实测约 30 cycles;常量内存是全局内存的只读区,走每 SM 8KB 专用缓存,它与共享内存同住那块 192KB 统一 L1/共享 SRAM,命中约 30 cycles 而非几个周期——便宜在 warp 读同址只发一次广播事务,省事务数不省延迟;全局内存约 400-500 cycles。延迟比 400÷30≈13×;带宽比:共享内存每 SM 每拍 128 Byte,108×128×1.41 GHz≈19.5 TB/s 片上,对 1.555 TB/s 片外,约 12.5×。作用域决定同步:只有同 block 的线程能靠共享内存加 __syncthreads() 交换数据。两个杀手:bank conflict——32 bank×4 Byte,同 warp 命中同 bank 不同址会串行化 n 次(同址是广播,不算冲突);coalescing——warp 内 32 线程读连续 128 Byte 只花一次事务,跨步最多退化成 32 次。出处:Hwu PMPP 第4版第4章;Hennessy & Patterson 第2章。

第二幕 · 第一崩塌(带宽墙)

Roofline · 算术强度

算得快不等于跑得快 · 拐点6.26 FLOP/Byte · 核心之伤形成

你站在带宽之壁面前。壁面比想象的更近、更高、更冷。诺依曼老祖按下"全量运行"——你的矩阵乘法、归约、排序内核同时启动。6912个核心亮起紫光。但数据从HBM涌向核心的通道——只有1550 GB/s。

带宽之壁 带宽之壁

「你的核心算力是9.7 TFLOPS——FP64向量峰值。你的带宽是1550 GB/s。算术强度拐点是9700÷1550=6.26 FLOP per Byte。你的矩阵乘法分了块,tile=16,算术强度16÷4=4。4 < 6.26。你的核心在等数据。你被钉在这里了。」

诺依曼老祖 诺依曼老祖

「你被钉住了。但墙上有刻度——Roofline。读懂它,才知道瓶颈在哪。」

Roofline 算术强度追踪器

拖tile滑块 → 矩阵乘法标记沿X轴移动 → 过拐点6.26裂缝愈合

A100: FP64 峰值 9.7 TFLOPS · HBM2e 1550 GB/s · 拐点 9700÷1550 = 6.26 FLOP/Byte
tile_size = 16
算术强度 = 4.00 FLOP/Byte (内存受限)
详解 · K23-08 矩阵乘法并行化:分块算法与GPU实现

C=A×B(n×n)是标准考题。工作量 2n³ FLOP(FMA 记 2)。朴素内核每线程算一个 C[i][j],读 A 行 B 列各 n 个 float:2n FLOP 换 8n Byte,AI=0.25,且与 n 无关——矩阵放大一万倍也救不了不复用的写法。完美复用下 A、B、C 各过一遍 DRAM,流量 12n²,AI=2n³÷12n²=n/6,n=1024 时约 171。分块赚的就是这段:一个 block 负责 T×T 输出块,每阶段把 A、B 各一个 T×T 子块搬进共享内存(8T² Byte)复用。按单块算账:工作 2T²n FLOP、流量 A 行条+B 列条=8Tn Byte,故 AI=T/4;系数 1/4 与拐点同口径(都按 FMA=2),按 tile/8 算就是分子少记一半 FLOP。代进 A100:tile=16→AI=4.00,实得 1550×4=6.2 TFLOPS 即峰值 64%,正是台上那道裂缝;tile=26→6.5 刚过 6.26;tile=64→16,被 9.7 封顶。T 也不能无限大:共享内存要 8T²,T=128 需 128KB,已贴着每 block 164KB 上限。出处:Hwu PMPP 第4版第4、5章;Hennessy & Patterson 第4章。

详解 · K23-15 GPU计算的性能瓶颈:带宽限制/分支分歧/合并访存

GPU 内核跑不满,原因无非三类。一是带宽:拐点左侧性能=带宽×AI,与核心数无关。台上三个内核代进 A100(拐点 6.26)——矩阵乘法 AI=4.00→6.2 TFLOPS(64%);归约 AI=0.25→388 GFLOPS(4%);排序 AI≈0.2→310 GFLOPS(3.2%)。二是分支分歧:warp 的 32 条 lane 走不同路径时,硬件按路径依次重放、用活跃掩码关掉无关 lane,两路 if-else 各占一半则有效算力只剩 50%——台上排序标的 67% 分歧率、33% 有效率即此意。Volta 后每线程有独立 PC,但性能上锁步仍是唯一快路。解法:重排数据让同 warp 走同一路径;用 min/max 或 __ballot_sync 谓词化;选比较交换位置与数据无关的算法(bitonic 网络编译期定死)。三是访存模式:合并访问与 bank conflict,机理见 K23-06,跨步访存最多把一次事务退化成 32 次。三类都不进 FLOP 计数——只看峰值算力采购必翻车,Nsight Compute 的 speed-of-light 才是裁判。出处:Hwu PMPP 第4版第4、5章。

第二幕 · 方法觉醒

分解之术 · 归约与排序

工作-跨度模型 T₁/T∞ · Brent定理 · 归约树最后warp串行

网格之海深处,一块巨大的石板上刻着一张DAG——节点是任务,边是依赖。诺依曼老祖指着DAG,带宽之壁的光在边上流动。

诺依曼老祖 诺依曼老祖

「并行不是多开几个核心跑循环。你得先把问题分解成没有依赖的子任务。这条最长路径,叫关键路径——它决定你最少要花多少时间。」

带宽之壁 带宽之壁

「归约的前几轮确实漂亮。但最后几轮呢?剩32个数时,一个warp就够了——6912个核心里只有32条lane在干活,另外6880个在等。」

DAG分解 · 归约树 · Brent定理

拆DAG从串行链变树形 → 拖p滑块验证Brent → 归约树最后warp串行

DAG 分解 · 串行链 → 树形

详解 · K23-07 并行算法设计:分解策略/工作-跨度模型

并行算法第一步不是写循环,是分解。三视角:任务分解(拆独立子任务)、数据分解(切数据域,最常见——CUDA grid-stride loop、MPI 块划分属此类)、流水分解(各级重叠)。拆完得一张 DAG,用工作-跨度模型量它:T₁=所有结点权重之和(work,单核跑完的时间),T∞=最长路径(span/critical path,无穷多核也压不下去的下限),并行度=T₁/T∞。Brent 定理(Richard P. Brent,1974,SIAM J. Comput.)夹住两头:T₁/p ≤ Tₚ ≤ T₁/p + T∞。右边才是定理本体,证明靠贪心调度器——每步让空闲处理器各取一个就绪结点:这步要么有 p 个结点在跑(满步),要么就绪集已空(闲步);满步≤T₁/p、闲步≤T∞,相加即上界。推论:p≤T₁/T∞ 时 Tₚ≤2T₁/p,至少拿到线性加速的一半;超过并行度后加核只加电费。台上例子:8 个数求和共 7 次加法,串行链 T∞=7;二叉归约树 8 叶加 7 内部结点,T₁ 仍是 7、T∞=log₂8=3、并行度 7/3≈2.33,p=3 就到顶。出处:Hennessy & Patterson 第4、5章;CLRS 第27章。

详解 · K23-09 并行归约(Reduction)算法

归约是把 n 个元素用满足结合律的算子合成一个标量:求和、最值、点积都是它。串行版 n-1 次运算、深度也是 n-1。二叉树归约把深度压到 ⌈log₂n⌉,运算次数一次不少还是 n-1——工作量守恒。台上 n=1024 手算:1023 次加法分 10 轮,各轮活跃线程 512、256、128、64、32、16、8、4、2、1;前 5 轮合计 992 次占 96.97%,后 5 轮合计 31 次占 3.03%,且第 6 轮起整条尾巴塞在一个 warp 里——6912 条 lane 只剩 32 条在动,末尾并行度 32/6912≈0.46%。平均并行度 T₁/T∞=1023/10≈102,离 6912 差 68 倍——这就是"归约吃不满 GPU"的定量说法。工程实现两级:block 内共享内存树形归约、每块写一个部分和回全局,再由第二个内核(或 atomicAdd)合并;最后 5 轮改用 __shfl_down_sync 寄存器间直传,省掉共享内存往返。还有个 HPC 老坑:浮点加法不满足结合律,(a+b)+c≠a+(b+c),改树形就改结果——要位级可复现须固定树形或用 Kahan 补偿求和。出处:Hwu PMPP 第4版归约与前缀和两章。

详解 · K23-10 并行排序算法:Bitonic Sort/Radix Sort

比较排序的 Ω(n log n) 下界是给串行算法的;并行要的是浅深度加规则访存。Bitonic sort 由 Ken Batcher 1968 年提出(AFIPS《Sorting Networks and Their Applications》):先把序列整成双调(先升后降)再折半归并。n 个元素需 log₂n 阶段、第 k 阶段 k 步、每步 n/2 个比较交换:n=1024 时 10 阶段、Σk=55 步、55×512=28,160 次,比最优 n·log₂n=10,240 多花 2.75 倍,换来的是比较交换位置与数据完全无关(oblivious sorting network)。这对 GPU 是决定性的:控制流编译期定死、无数据依赖分支,可整片展开,剩下的分歧只在"换不换",用 min/max 即可谓词化。Radix sort 不比较,按基数逐位分桶:32 位整数每趟 8 位分 256 桶、4 趟搞定;每趟是一次计数排序——直方图、前缀和、scatter,单趟 O(n+k)、总 O(kn),前缀和用 Blelloch 并行 scan(深度 O(log n))。CUB/Thrust 的 DeviceRadixSort 是 GPU 默认选择,bitonic 更适合小 n 或需要固定网络的场合。出处:Hennessy & Patterson 第4章;Batcher 1968。

第二幕 · 范式之辨

双范式 · OpenMP 与 MPI

共享内存 vs 分布式内存 · #pragma omp vs MPI_Send/Recv

矩阵墙后是一条分岔路——左边通向一台多核机器,右边通向一排用光缆相连的节点。诺依曼老祖递来两把钥匙。

诺依曼老祖 诺依曼老祖

「OpenMP管一台机器里的核心——共享内存,线程同看一份数据。MPI管多台机器之间——分布式内存,每个进程有独立地址空间,要显式传消息。」

带宽之壁 带宽之壁

「共享内存有它的墙——核多了要锁。分布式内存也有它的墙——节点间通信就是另一道带宽之壁。」

OpenMP vs MPI · 范式对照

切换范式 → 同一段代码两种写法 → 共享 vs 分布式特征卡归位

详解 · K23-11 并行编程模型:共享内存(OpenMP)/消息传递(MPI)

两种范式对应两种地址空间。OpenMP 走共享内存:1997 年 ARB 发布 1.0(先 Fortran,次年补 C/C++),当前 5.2(2021);fork-join 模型,主线程遇 #pragma omp parallel 派生线程、区末汇合。子句:for 分迭代、task 建任务图、reduction(+:sum)(每线程私有副本、出口归并,比 critical/atomic 快)、schedule(static|dynamic);线程数由 OMP_NUM_THREADS 定。切记 reduction 须与 #pragma 同行——预处理指令到行尾即止,跨行要用反斜杠续行。MPI 走分布式内存:Forum 1992 年成立,MPI-1.0 于 1994 年发布,MPI-2(1997)加单边通信、MPI-3(2012)加非阻塞集合通信、MPI-4.0 在 2021 年。点对点 MPI_Send/MPI_Recv 六参数(buf、count、datatype、dest/src、tag、comm),集合通信 Allreduce/Scatter/Gather,结尾必须 MPI_Finalize。通信用 α-β 模型:T=α+β·n(α 启动延迟、β 每字节传输时间)——所以小消息要合并、AllReduce 走 ring(通信量 2(p-1)/p·n 带宽最优,2017 年由 Horovod 带进深度学习)。TOP500 标配混合:节点内 OpenMP 或 CUDA、节点间 MPI。出处:Hennessy & Patterson 第5、6章;Gropp/Lusk/Skene《Using MPI》。

第三幕 · 第二崩塌 · 理性觉醒

性能之秤 · Amdahl定律

S=1/(f+(1-f)/p) · 串行比例f决定上限 · 核心之伤修复

带宽之壁的表面浮出一杆秤——一端是计算,一端是访存。诺依曼老祖站在秤边,带宽之壁的声音从壁面渗出。

诺依曼老祖 诺依曼老祖

「你的内核有串行段——初始化、最终归并、同步屏障。Amdahl说:如果10%的代码是串行的,不管你有多少核心,最多只能快10倍。」

带宽之壁 带宽之壁

「你补不补那道裂缝?裂缝在问你——你的矩阵乘法,算术强度到底卡在哪里?」

Amdahl 定律推导组装器

拖公式块组装5步推导链 → 拖f滑块看曲线 → 补核心之伤

详解 · K23-13 Amdahl定律在并行系统中的深入应用

Amdahl 定律出自 Gene Amdahl 1967 年 AFIPS 春季会议的《Validity of the Single Processor Approach to Achieving Large Scale Computing Capability》,本意是给"造一台超强单处理器"泼冷水。设串行比例 f、核数 p:S(p)=1/(f+(1-f)/p),p→∞ 得 S_max=1/f。手算:f=50%→2、10%→10、5%→20、1%→100、0.1%→1000。效率 E=S/p 掉得更快:f=10% 时 p=100 得 S=1/(0.1+0.009)=9.17、E=9.2%——百倍核心只换 9 倍性能。1988 年 Gustafson 在 CACM 发《Reevaluating Amdahl's Law》反驳:Amdahl 假设问题规模固定,现实却是先有更大的机器再解更大的问题;他把尺度反过来得 S(p)=p-f_s(p-1),f_s 是在 p 台处理器上按放大后的问题实测的串行时间占比。Sandia 数据:p=1024 时串行占比不到 1%、实测加速 1016-1021,代 f_s=0.004 得 1024-0.004×1023≈1019.9,自洽。口径提醒:两个 f 不是一个量,台上同一个滑块喂两条公式只是教学近似。第三堵墙两个公式都没写:Roofline——它们只分配时间份额,不管每核拿到峰值的几成;AI 低于拐点时性能钉在带宽×AI,弱扩展也逃不出。出处:Hennessy & Patterson 第1章;Gustafson, CACM 31(5), 1988。

第三幕 · 最高潮

Boss 战 · 带宽之壁三问

三问对决 · 排除刻痕收网 · 雾账清算 · 算术强度是灵魂

带宽之壁从远处的悬崖向你压来,凝成一道人形光壁——通体钢灰,没有五官,只有数据流在表面撞击碎裂的纹路。它开口了,声音像所有数据撞墙的回声。

带宽之壁 带宽之壁

「我不住在计算慢里。我住在你'算得快就够了'的幻觉里。你在每一处都欠了我的账——现在,一笔笔还。」

渊账累计 0 笔 · 三问全对后清零

第三幕 · 收束

通天之境 · HPC远眺

集群拓扑 · Slingshot/InfiniBand互连 · 量子涟漪 · 化神一层·并行通天·圆满

壁面碎成金色尘埃,视野豁然开朗——你看见无数计算节点以光缆相连,排列成一座通天之塔。诺依曼老祖递来两把钥匙,远处一道量子涟漪在空间中一闪而过。

诺依曼老祖 诺依曼老祖

「单机GPU之上,还有HPC集群。Frontier超算,9408个节点,Rmax 1.194 EFLOP/s——2022年首破E级、蝉联TOP500榜首两年,直到2024年6月被El Capitan的1.742 EFLOP/s接过去。但它的每个节点之间,靠Slingshot-11高速网络传消息,每端口200 Gb/s——那又是另一道带宽之壁。」

HPC集群拓扑 · 知识图谱点亮

拼集群拓扑 → 点亮知识碎片 → 境界突破

详解 · K23-14 高性能计算(HPC)集群与超级计算机架构

一台超算就是台上那张拓扑:计算节点(一颗 host CPU 加若干加速器)→ 高速互连 → 并行文件系统(Lustre)→ 作业调度器(Slurm/PBS)。榜单口径:TOP500 自 1993 年起每年 6 月、11 月各发一次,按 HPL(High-Performance Linpack)的 Rmax 排名。里程碑:1997 年 ASCI Red 首破 1 TFLOP/s,2008 年 Roadrunner 首破 1 PFLOP/s(1.026),2013-2016 天河二号 33.86 PFLOP/s 蝉联六届,2016-2018 神威·太湖之光 93.01 PFLOP/s(全国产申威 SW26010),2022 年 6 月 Frontier(ORNL,HPE Cray EX,9408 节点、每节点 1 颗 AMD EPYC 加 4 张 MI250X)以 1.102 EFLOP/s 成为首台公开确认的 E 级机、同年 11 月刷到 1.194,2024 年 6 月 El Capitan(LLNL,AMD MI300A APU)以 1.742 EFLOP/s 接棒——所以"Frontier 全球第一"必须带时间戳。互连是集群的第二道带宽之壁:Frontier 用 HPE Slingshot-11(每端口 200 Gb/s、每节点 8 张网卡),不是 InfiniBand;InfiniBand NDR 已到 400 Gb/s(Summit、Leonardo 用它),Fugaku 走自研 Tofu-D 六维环面。能效另有 Green500:Frontier 曾以 62.68 GFLOP/s/W 于 2022 年 11 月同时登顶两榜,整机 22.7 MW。出处:Hennessy & Patterson 第6章;TOP500.org。

复盘 · 并行通天主线

S0 下注:加速比全由算术强度定,AI=0.083 的向量加法只拿峰值 1.3%,理想矩阵乘法 AI≈171 跑满(K23-12)。S1 频率墙:P=α·C·V²·f,V 随 f 涨即 P∝f³,3.5 GHz/100 W 推到 8 GHz 是 1194 W;Dennard 缩放 1974 年生、2006 年卒,Flynn 1966 年四象限是出路地图(K23-01/02)。S2 SIMT:warp=32 lane 锁步,6912 是 lane 数不是 6912 个 CPU;16 warp/调度器盖不住 400 cycles 的 DRAM 延迟(K23-03/04)。S3 CUDA:全局索引 threadIdx.x+blockIdx.x*blockDim.x,block 须 32 的倍数;1024 线程只占满 15% 核心(K23-05)。S4 内存塔:1/30/30/400 cycles,延迟差 13×、带宽差约 12.5×(K23-06)。S5 Roofline:拐点 9.7÷1.55=6.26,分块 AI=tile/4,tile=16→4.00 得峰值 64%,tile≥26 才过拐点(K23-08/15)。S6 分解:Brent 夹住 T₁/p ≤ Tₚ ≤ T₁/p+T∞;8 个数求和 T₁=7、T∞ 从 7 压到 3;1024 元素归约前 5 轮吃掉 97%,尾巴 31 次全在 1 个 warp 里;bitonic n=1024 要 28,160 次比较交换换数据无关的网络(K23-07/09/10)。S7 两范式:OpenMP 靠 reduction 子句,MPI 按 T=α+β·n 收通信税(K23-11)。S8 两堵墙:Amdahl f=10% 时 p=100 也只到 9.17;Gustafson S=p-f(p-1) 靠弱扩展绕开,但两个 f 不是一个量;Roofline 那堵两个公式都没写(K23-13)。S10 远眺:Frontier 9408 节点、Slingshot-11、1.194 EFLOP/s(2022-2023 榜首),2024 年 6 月被 El Capitan 的 1.742 接棒(K23-14)。一句话:核心可以堆,带宽不能堆——并行的灵魂是每 Byte 访存能喂多少次运算。

交互实验

并行加速:Amdahl vs Gustafson

对比阿姆达尔(固定问题)与古斯塔夫森(固定时间)两种加速视角。