TPU v1 完全拆解(二):脉动阵列、weight-stationary 数据流与 Roofline 性能归因

9276 字
46 分钟
TPU v1 完全拆解(二):脉动阵列、weight-stationary 数据流与 Roofline 性能归因

AI 生成内容声明

系列回顾与本文路线图#

上一篇从整体架构讲清了 TPU v1(2015 年部署进 Google 数据中心的推理专用 ASIC,2017 年 ISCA 论文公开)的三件事:为什么推理需要一颗”只做矩阵乘”的专用芯片、芯片由哪些模块组成(256×256 的 MXU、24 MiB Unified Buffer、8 GiB DDR3 Weight Memory)、以及 CISC 风格的指令集如何用一条指令算完一层网络。那篇留下了三个悬而未决的问题:

  1. 92 TOPS 是怎么算出来的——“256×256 阵列 × 700 MHz”背后,阵列内部到底是什么样的计算结构,数据怎么流动?
  2. 为什么平均只用 28%——论文实测 6 个生产模型里,MXU 平均只有 28% 的时钟周期在真正做矩阵运算(CNN 能到 78%,MLP 只有 10% 左右),这个巨大的落差从哪来?
  3. “带宽限制”到底怎么定量——论文反复说 MLP/LSTM 是”memory-bound”、CNN 是”compute-bound”,这句话的数学依据是什么?

这三个问题恰好是同一枚硬币的两面:脉动阵列(systolic array)决定了 MXU 内部怎么算,Roofline 模型决定了这个”怎么算”能跑多快。本文就把这两套机制彻底拆开:先从 1978 年 Kung 和 Leiserson 提出脉动阵列的动机讲起,再到 TPU 的 weight-stationary 数据流时序,然后用 Roofline 模型完成对”28% 利用率”的完整性能归因,最后用论文的敏感性实验回答”如果 Google 多给 TPU 六个月,它会变成什么样”。

背景:为什么要发明脉动阵列#

数据移动比计算更贵#

脉动阵列不是 TPU 的发明,它 1978 年就由卡内基梅隆大学的 H. T. Kung 和 C. E. Leiserson 提出(Kung 在 1982 年的《Why Systolic Architectures?》中做了系统总结)。要理解它为什么出现,先要理解当时的困境,而这个困境到现在依然是 AI 芯片设计的核心约束。

传统处理器(冯·诺依曼结构)执行一条乘加指令的流程是:取指令 → 从寄存器/内存取操作数 → 计算 → 写回结果。对矩阵乘来说,这个流程有个致命问题:数据搬运的时间和能量远大于计算本身。TPU 论文里引用了一个关键事实:读一次大型 SRAM 消耗的能量,比做一次算术运算还要多(论文引用 [Kun80][Ram91][Ovt15b] 三个出处支持这个论断)。具体量级可以参考后来的经典数据(Horowitz 2014 年发表在 ISSCC 的功耗表):一次 8 位整型乘法约 0.2 pJ,而一次 32 KB SRAM 读取约 5 pJ,是乘法的 25 倍;DRAM 读取更是 640 pJ,是乘法的 3200 倍。

Kung 的洞察是:如果能减少”把数据搬来搬去”的次数,即使计算单元简单一点,整体能耗和性能也会大幅改善。他注意到很多信号处理应用(FIR 滤波、矩阵乘、卷积)有一个共同特点——数据复用率极高:同一个输入要参与很多次不同的乘加。如果能把数据”留在”计算单元附近反复使用,而不是每次用完就存回内存、下次再取,就可以把数据搬运量缩小一个数量级。

“心脏泵血”式的解决方案#

Kung 给出的方案是:把大量简单的乘加单元(PE,Processing Element)排列成规则网格,让数据像血液一样从阵列边缘流入,流经所有 PE 后从另一侧流出,中间不停顿、不写回内存。这就是”systolic”(脉动的)这个名字的由来——数据在阵列中一波一波地推进,如同心脏泵血。

脉动阵列相对传统结构有三个根本区别:

  1. 数据在 PE 之间直接传递,不走内存。每个 PE 只和相邻的 PE 通信(连线短、功耗低、布线简单),计算结果通过相邻 PE 接力传递,只有最终结果才写回。
  2. 每个 PE 极其简单。一个 PE 只需要一个乘法器、一个加法器/累加器、几个寄存器,没有指令、没有控制逻辑。整个阵列用同一个控制信号锁步(lockstep)运行,所有 PE 在同一拍做同一件事。SIMD 控制方式下,控制逻辑由整个阵列共享,平均到每个 PE 上的开销几乎可以忽略——对比 MIMD 架构(每个 PE 自带一套指令解码和调度逻辑,开销是数千晶体管量级),省下的面积全部可以换成计算单元。
  3. 数据复用发生在 PE 内部。一个数被读进阵列后,在流动过程中被沿途的所有 PE 反复使用,不需要重复访问寄存器堆或内存。

用论文里的说法:TPU 的 MXU 使用脉动执行(systolic execution)来省电,手段是减少对 Unified Buffer 的读写次数——UB 是 24 MiB 的大 SRAM,读写它比做算术贵得多。脉动阵列让一个激活值被读入一次后,在阵列里被 256 个 PE 轮流使用,把”读一次数据做一次运算”变成”读一次数据做几百次运算”。

顺带说一个历史背景:脉动阵列在 20 世纪 80–90 年代曾被 CMU 的 Warp、iWarp 等项目做成真实芯片,但最终没有商业成功——原因不是硬件不行,而是编程太难:通用 CPU 只需要程序员写串行代码,而脉动阵列要求程序员/编译器把算法显式映射成数据流动的时序,当时没有编译器能做好这件事。Kung 本人也在论文里承认这一点。TPU 之所以能让脉动阵列”复活”,恰恰是因为它把软件问题绕开了:推理模型被固定成”权重静态、数据流式”这一种模式,编译器只需要处理矩阵乘一种映射,而矩阵乘到脉动阵列的映射是数学上最规整的。领域专用设计的另一个侧面是:它把”难的部分”从硬件挪到了软件栈,而软件栈只为一类负载服务

脉动阵列的三种数据流模式#

脉动阵列的 PE 网络有了,但具体”让哪个数据流动、哪个数据驻留”,有好几种设计选择。学术界按”驻留在 PE 里的数据是什么”把数据流(dataflow)分成三种经典模式。这里先用一张社区解读图把三种模式放在一起对比(图片来源:博客园《systolic array 架构设计综述》):

脉动阵列三种数据流模式对比:input-stationary(输入驻留)、weight-stationary(权重驻留)、output-stationary(输出驻留),黄色表示驻留在 PE 内的数据(图片来源:博客园《systolic array 架构设计综述》)
脉动阵列三种数据流模式对比:input-stationary(输入驻留)、weight-stationary(权重驻留)、output-stationary(输出驻留),黄色表示驻留在 PE 内的数据(图片来源:博客园《systolic array 架构设计综述》)

三种模式各自的取舍:

Input-Stationary(输入驻留,IS):输入激活值(activation)存在 PE 里不动,权重在阵列中流动。适合激活值复用率极高的场景(比如卷积的输入特征图被多个卷积核共享),但权重流动意味着每个 PE 都要接收新的权重,对权重带宽的要求高。

Weight-Stationary(权重驻留,WS):权重预先加载进 PE 的寄存器,之后一直驻留不动;激活值和部分和在阵列中流动。每个权重被读进 PE 一次,之后被反复使用。这是 TPU 的选择。适合权重复用率高的场景——恰好就是推理:推理时权重是静态的、只读的,而且同一批权重要被一批(batch)样本反复使用。

Output-Stationary(输出驻留,OS):每个 PE 负责一个输出元素,部分和一直驻留在 PE 里累加,直到最终结果算完才写回。适合输出数据复用率高的场景(部分和反复累加),但每个 PE 要保存完整的部分和,寄存器开销较大。

Eyeriss(MIT 2016 年 ISCA 发表的著名 CNN 加速器,论文系统性比较了这三种数据流)给出了一个定量结论:不存在”对所有网络都最优”的数据流——WS 适合权重复用高的层,OS 适合输出累加多的层,Eyeriss 为此提出了一种自适应混合的 row-stationary 数据流,在 AlexNet 的卷积层上比单一数据流节能 1.4–2.5 倍。

为什么 TPU 选 weight-stationary#

TPU 选 WS 的理由非常直白,结合它的目标负载(推理)可以逐条列出:

  1. 推理的权重是只读的。推理不更新权重,权重从 DDR3 读进阵列后永远不需要写回,天然适合”驻留”。
  2. 权重复用次数 = 批大小。一层全连接层,权重矩阵是 N×NN \times N(8 位 = N×NN \times N 字节),一批 BB 个样本的矩阵乘是 B×N×NB \times N \times N 次乘加。权重驻留时,每个权重字节被复用 BB,于是”每读取 1 字节权重能换来 BB 次运算”。论文 Table 1 给出的操作强度(ops/weight byte)恰好等于各模型的批大小:MLP0 是 200、MLP1 是 168、LSTM0 是 64、LSTM1 是 96——和批大小完全一致,这不是巧合,这正是 WS 数据流的定义性质。
  3. 激活值沿行流动,天然被复用 256 次。一个激活值从左侧进入阵列后,会在这一行经过 256 个 PE(对应 256 个不同的权重),相当于一次 UB 读取支撑 256 次乘加。UB 读写的次数被压缩了约 256 倍——这正是论文说的”用脉动执行减少对 UB 的读写来省电”。

对比之下,如果选 OS 模式,部分和驻留在 PE 里,激活值和权重都要流动,对权重带宽的要求就上去了;如果选 IS 模式,权重流动会破坏”权重只读一次”的最大优势。在”静态权重 + 大批量推理”这个特定场景里,WS 是数学上最优的选择。

MXU 的 weight-stationary 数据流:时序原理#

论文的 Figure 4 给出了 MXU 内部的数据流图,这是理解整颗芯片心脏的关键图(图片来源:TPU 论文 Figure 4):

TPU 论文 Figure 4:MXU 的脉动数据流——激活从左侧流入、权重从顶部加载、部分和向下流动(图片来源:TPU 论文 Figure 4)
TPU 论文 Figure 4:MXU 的脉动数据流——激活从左侧流入、权重从顶部加载、部分和向下流动(图片来源:TPU 论文 Figure 4)

论文正文对这张图的描述值得逐句拆解:

“数据从左侧流入,权重从顶部加载。一个 256 元素的乘加运算以对角波前(diagonal wavefront)的方式穿过整个矩阵。权重预先加载,并随着新数据块的到来与波前同步生效。控制与数据被流水化,从而制造出’256 个输入同时读出、并瞬间更新 256 个累加器中的一个位置’的假象。”

把这段话翻译成时序逻辑,MXU 的一次矩阵乘指令(B×256 激活 × 256×256 权重 → B×256 部分和)分三个阶段:

阶段一:权重预加载(Weight Preload)#

MXU 内部能放 64 KiB 的权重块(tile,恰好是 256×256 个 8 位权重)。执行矩阵乘之前,权重必须先搬进阵列——从 Weight FIFO 逐拍移入,256 字节/拍,正好 256 拍移完一个 tile(256×256 字节 ÷ 256 字节/拍)。这 256 拍的加载时间就是论文里说的”把 tile 移进来需要 256 拍”(the 256 cycles it takes to shift a tile in)。

为了掩盖这 256 拍,MXU 里放了两份权重 tile 做双缓冲:一个 tile 在计算时,另一个 tile 在后台从 FIFO 加载。计算一个 tile 需要 B 拍(见阶段三),只要 B>256B > 256,双缓冲就能完全隐藏权重加载时间。论文选 4096 组累加器的推导也与此相关:达到峰值需要的操作强度约 1350,凑整到 2048,再翻倍到 4096,就是为了让编译器在跑峰值时还有空间做双缓冲。

阶段二:波前传播(Wavefront Propagation)#

权重就位后,激活数据开始从左侧流入。但注意:流进阵列的激活不是立刻全部生效的。左上角的 PE 第一拍就拿到了第一个激活和第一个权重,但第 256 列的 PE 要等激活沿着行传播 255 拍才能拿到数据。于是每个”计算前沿”在阵列里呈对角线形状推进——这就是”对角波前”。一个 256 元素乘加操作的波前从左上角到右下角需要约 512 拍(256 拍横移 + 256 拍纵移)的流水填充(pipeline fill)时间,之后阵列进入稳态,每拍吞吐 256 个乘加。

论文强调”控制与数据被流水化,制造出 256 个输入同时读出的假象”——意思是:从软件/编译器视角看,MXU 每拍读入 256 个激活、瞬间产生 256 个部分和,内部的波前传播延迟被完全隐藏了。正确性层面,软件根本感知不到脉动结构的存在;但性能层面,软件要操心阵列的延迟(比如层与层之间的依赖、权重加载的节奏)。

这里值得停下来和 GPU 的执行模型做一次对照,因为两者的本质差异正好解释了后续所有性能数据。GPU 的 SIMT 模型里,一个矩阵乘被拆成几十万个线程,每个线程独立地从寄存器/显存取数、计算、写回——并行度来自大量独立线程,控制分散在每个线程里;而脉动阵列的并行度来自数据在空间上的流动:65536 个 PE 不需要知道自己在算矩阵的哪一块,它们只是接收邻居的数据、做乘加、传给下一个邻居。GPU 靠”很多简单的执行单元 + 复杂调度”赢下通用性,脉动阵列靠”规则数据流 + 几乎为零的调度”赢下面积和能耗。代价是:脉动阵列的延迟不是可调度的,而是几何的——波前要横穿整个阵列(约 512 拍),这条延迟路径无法像 GPU 那样用更多线程来隐藏,只能靠编译器预先排好数据到达的节奏。

阶段三:稳态流水(Steady State)#

每拍,MXU 从左侧读入 256 个激活值,产生 256 个部分和,写入下方 4 MiB 累加器中的一组。B×256 的输出需要 B 拍算完——论文说的”MatrixMultiply 指令耗时 B 个周期”就是这个意思。之所以是 B 而不是 B×256,是因为 256 列的计算是并行的。计算期间权重不需要逐拍读取——它们已经驻留在 PE 里了,后台同时以 256 字节/拍的速率把下一个权重 tile 移入双缓冲,为下一个矩阵乘做准备。

峰值算力的来源此时一目了然——每拍完成 256×256 次乘加,乘上时钟频率:

Ppeak=256×256×fclk=65536×700 MHz=45.9 TeraMAC/s=92 TOPSP_{peak} = 256 \times 256 \times f_{clk} = 65536 \times 700 \text{ MHz} = 45.9 \text{ TeraMAC/s} = 92 \text{ TOPS}

注意这里的换算:65,536 个 MAC 单元每拍做 65,536 次乘加,每个 MAC 记 2 次运算(乘 + 加),所以 45.9 TeraMAC/s = 91.8 ≈ 92 TOPS。这个”每个 MAC 算 2 次运算”的计数口径在后面 Roofline 分析里还会回来找我们麻烦,先记住。

一个关键问题:为什么 B 拍算完一层,阵列还是会闲着?#

严格说,上述时序成立的前提是权重永远跟得上。权重从 DDR3 到 Weight FIFO 再到 MXU,路径上的带宽是 34 GB/s(DDR3);而 MXU 消费权重的速率是多少?每个 tile 是 64 KiB,B 拍算完,所以权重的消费速率是:

64 KiBB/700 MHz=65536 BB×1.43 ns\frac{64 \text{ KiB}}{B / 700 \text{ MHz}} = \frac{65536 \text{ B}}{B \times 1.43 \text{ ns}}

B=200B = 200(MLP0 的批大小)时,消费速率约 229 GB/s——是 DDR3 供给带宽(34 GB/s)的 6.7 倍。权重供给严重不足,MXU 必然频繁空转等待。Weight FIFO 只有 4 个 tile 深(4×64 KiB),最多预取 4 个 tile,DRAM 延迟和带宽缺口全靠这 256 KiB 的缓冲来扛,扛不住的部分就变成 stall。这就是”memory-bound”的微观图景,下一节用 Roofline 把它变成精确的数学。

Roofline:判断”卡算力还是卡带宽”的数学工具#

模型定义#

Roofline 模型由 Berkeley 的 Williams、Waterman 和 Patterson 在 2009 年提出(原论文《Roofline: An Insightful Visual Performance Model for Multicore Architectures》),是性能归因的通用工具。它只基于两个硬件参数画图:

  • 峰值算力 PpeakP_{peak}(纵轴天花板,水平线):硬件每秒最多能做多少次运算;
  • 内存带宽 BmemB_{mem}(斜线):硬件每秒最多能从内存取多少字节。

横轴是操作强度(operational intensity)II,定义为”每取 1 字节数据能做多少次运算”。实际可达到的性能被两者夹逼:

P=min(Ppeak, I×Bmem)P = \min(P_{peak},\ I \times B_{mem})

I×Bmem<PpeakI \times B_{mem} < P_{peak} 时,性能被带宽卡住(斜线区,memory-bound);反之被算力卡住(水平区,compute-bound)。两条线的交点叫脊点(ridge point):

Iridge=PpeakBmemI_{ridge} = \frac{P_{peak}}{B_{mem}}

脊点左侧的程序提高算力没用,只能提高数据复用率(增大 II)或带宽;脊点右侧的程序提高带宽没用,只能提高算力。Roofline 的最大价值是告诉你”优化该往哪个方向使劲”

为推理定制的 Roofline:操作强度重定义#

TPU 论文对 Roofline 做了一处关键改造:标准 Roofline 的强度分母是”从内存取的所有数据”,而推理场景下权重是只读的、激活和中间结果都存在片上 24 MiB UB 里,数据搬运的大头只有权重。所以论文把操作强度重新定义为:

I=总整数运算量(MAC 数)从 Weight Memory 读取的权重字节数I = \frac{\text{总整数运算量(MAC 数)}}{\text{从 Weight Memory 读取的权重字节数}}

这个定义下,对一层全连接层,I=BI = B(每字节权重被 batch 个样本复用);对卷积层还要乘上空间复用因子,所以 CNN 的强度远高于 MLP。论文 Table 1 给出的六个模型强度:MLP0 = 200、MLP1 = 168、LSTM0 = 64、LSTM1 = 96、CNN0 = 2888、CNN1 = 1750。

TPU 的 Roofline 图(论文 Figure 5,即下面的图)就是在 log-log 坐标系里画:水平线在 92 TOPS 处,斜线从原点出发,脊点远在右侧 1350 ops/byte 处:

TPU 论文 Figure 5:单颗 TPU die 的 Roofline 模型。脊点位于 1350 ops/weight-byte(右侧远处),图上有很长一段带宽限制的斜线区,六个模型的星点分布其中(图片来源:TPU 论文 Figure 5)
TPU 论文 Figure 5:单颗 TPU die 的 Roofline 模型。脊点位于 1350 ops/weight-byte(右侧远处),图上有很长一段带宽限制的斜线区,六个模型的星点分布其中(图片来源:TPU 论文 Figure 5)

论文原话:“TPU 的 Roofline 有一段很长的斜线部分,这意味着操作强度不高的应用性能受内存带宽限制而非峰值算力限制。六个应用中有五个正好顶着天花板:MLP 和 LSTM 是带宽受限(memory-bound),CNN 是算力受限(compute-bound)”

用数字验证一遍(顺便验证我们的强度定义是否自洽)。以 MLP0 为例:I=200I = 200 MAC/byte,Bmem=34B_{mem} = 34 GB/s,则带宽天花板为:

Pbound=I×Bmem=200×34 GB/s=6.8 TeraMAC/s=13.6 TOPSP_{bound} = I \times B_{mem} = 200 \times 34 \text{ GB/s} = 6.8 \text{ TeraMAC/s} = 13.6 \text{ TOPS}

论文实测 MLP0 是 12.3 TOPS——贴着 13.6 TOPS 的天花板跑,剩余差距来自流水线 stall 和激活搬运。CNN0:I=2888>1350I = 2888 > 1350,落在水平区,天花板就是 92 TOPS,实测 86 TOPS(93% 的峰值)。每一个数字都对得上,这就是 Roofline 作为分析工具的威力——不是定性描述,是精确到一位小数的定量预测。

这里顺带解开一个容易让人困惑的细节:按 Iridge=Ppeak/Bmem=92 TOPS/34 GB/sI_{ridge} = P_{peak} / B_{mem} = 92 \text{ TOPS} / 34 \text{ GB/s} 算,脊点应该是 2705 ops/byte,为什么论文图里标 1350?因为 92 TOPS 是”每个 MAC 计 2 次运算”的宣传口径,而 Roofline 图里的”运算”按 MAC 计数:46 TeraMAC/s ÷ 34 GB/s ≈ 1353 MAC/byte ≈ 1350。论文在 Table 1 里给的操作强度(MLP0 = 200 等)同样按 MAC 计数。理解这个口径差异很重要——今天各大厂商宣传的 TOPS 数字普遍按”乘+加=2 次运算”计算,比较芯片时稍不留神就会把口径不同的数字直接相除。

三芯片 Roofline 对比:为什么 TPU 赢在”别人够不到天花板”#

论文把三颗芯片的 Roofline 画在同一张 log-log 图上(论文 Figure 8):

论文 Figure 8:TPU(星)、K80 GPU(三角)、Haswell CPU(圆圈)的 Roofline 合并图。所有 TPU 的星点都位于另外两条 Roofline 之上(图片来源:TPU 论文 Figure 8)
论文 Figure 8:TPU(星)、K80 GPU(三角)、Haswell CPU(圆圈)的 Roofline 合并图。所有 TPU 的星点都位于另外两条 Roofline 之上(图片来源:TPU 论文 Figure 8)

三颗芯片的参数(论文 Table 2):

芯片峰值算力内存带宽脊点(ops/weight byte)
Intel Haswell E5-2699 v32.6 TOPS(FP)~200 GB/s13
NVIDIA K80(单 die)2.8 TOPS(8-bit)160 GB/s 有效9
TPU92 TOPS34 GB/s1350

K80 的 160 GB/s 是关闭 Boost、开启 SECDED 校验后的有效带宽(标称 240 GB/s);Haswell 的脊点 13 ops/byte 意味着它的应用只要强度超过 13 就落入算力限制区——而 TPU 的应用强度动辄 200+。换句话说:

  • CPU/GPU 的 Roofline 形状”矮而窄”:算力低、带宽高,脊点靠左,但应用实际跑不到自己的屋顶——被延迟约束压着(见下节);
  • TPU 的 Roofline”高而宽”:算力高、带宽低,脊点靠右,MLP/LSTM 这类低强度应用恰好落在长斜线区,虽然达不到峰值,但斜线区的高度本身就比 CPU/GPU 的整个屋顶还高

这就是论文那句”所有 TPU 的星点都位于其他两条 Roofline 之上”的含义——TPU 哪怕只发挥三成,也够碾压对手的巅峰状态。

延迟约束:CPU/GPU 为什么连自己的屋顶都够不到#

Roofline 给出的是”无约束下的性能上限”,但推理服务还有一个额外的约束——99 百分位响应时间。论文 Table 4 用 MLP0(要求 7 ms 内返回)展示了这个约束的杀伤力:

平台批大小99% 响应时间吞吐(IPS/die)相对最高吞吐
CPU167.2 ms5,48242%
CPU6421.3 ms13,194100%
GPU166.7 ms13,46137%
GPU648.3 ms36,465100%
TPU2007.0 ms225,00080%
TPU25010.0 ms280,000100%

批大小翻 4 倍,CPU/GPU 吞吐也翻 4 倍——理论上它们都能跑很高,但大批次把延迟顶爆了,应用被迫退回 batch=16,只发挥 37%–42%。TPU 因为单条指令就是一层网络、执行时间确定,敢用 batch=200 还压住延迟在 7.0 ms,发挥 80%。注意 GPU/TPU 的”100%“也并非自己的 Roofline 天花板——论文 Table 4 注明,GPU 和 TPU 的最大 MLP0 吞吐受主机服务器开销限制,它们离屋顶还差着一大截。CPU/GPU 输在两条线上:峰值不够高,且连自己的峰值都用不满。

性能归因:92 TOPS 的阵列为什么平均只用 28%#

Table 3:性能计数器的完整账本#

现在回答第二个问题。论文用硬件性能计数器(performance counter)给每个模型记了一本”阵列时间账”,这是整个性能归因的核心数据(论文 Table 3,行百分比为矩阵单元的时间分配或 MAC 利用率):

计数器(占矩阵单元周期比例)MLP0MLP1LSTM0LSTM1CNN0CNN1平均
阵列活跃周期(在做矩阵运算)12.7%10.6%8.2%10.5%78.2%46.2%28%
64K 阵列中有效 MAC 占比(% 峰值)12.5%9.4%8.2%6.3%78.2%22.5%23%
无效(空转)MAC0.3%1.2%0.0%4.2%0.0%23.7%5%
权重加载等待(weight stall)53.9%44.2%58.1%62.1%0.0%28.1%43%
权重移入阵列(weight shift)15.9%13.4%15.8%17.1%0.0%7.0%12%
非矩阵周期17.5%31.9%17.9%10.3%21.8%18.7%20%
RAW 流水线冒险停顿3.3%8.4%14.6%10.6%3.5%22.8%11%
输入数据停顿(PCIe)6.1%8.8%5.1%2.4%3.4%0.6%4%
实测性能(TOPS,峰值 92)12.39.73.72.886.014.121.4

第一行”阵列活跃周期”就是那个著名的 28% 平均利用率。按论文的说明,第一、四、五、六行(阵列活跃、权重加载等待、权重移入、非矩阵周期)是对矩阵单元周期活动的完整划分,四行加起来恒等于 100%(以 MLP0 验证:12.7 + 53.9 + 15.9 + 17.5 = 100.0);其余各行是更细的分解或原因解释,可以重叠。这张表的解读逻辑是:阵列闲着的时间,要么在等权重(stall + shift),要么在处理非矩阵事务(激活、同步),要么在等依赖(RAW)或等输入(PCIe)。下面按三类负载分别讲。

MLP/LSTM:权重带宽卡死的典型 memory-bound#

以 MLP0 为例(占当时部署 TPU 的 61%,它的权重就是平均值的主要贡献者):阵列只有 12.7% 的时间在算,53.9% 的时间在等权重从 DDR3 搬过来,15.9% 在把新权重块移入阵列——70% 的周期都耗在”喂权重”上。

为什么会等?回到上一节的数值:MLP0 的强度 I=200I = 200,带宽天花板 13.6 TOPS,而阵列本身能跑 92 TOPS。阵列的”饭量”是权重组装线的 6.7 倍。具体到一个 64 KiB 权重 tile:

  • 从 DDR3 搬一个 tile:65536 B/34 GB/s1.93 μs65536 \text{ B} / 34 \text{ GB/s} \approx 1.93\ \mu\text{s}
  • 阵列算完一个 tile(batch=200):200 拍/700 MHz0.29 μs200 \text{ 拍} / 700 \text{ MHz} \approx 0.29\ \mu\text{s}

搬进来要 1.93 µs,算掉只要 0.29 µs,搬运时间是计算时间的 6.7 倍。Weight FIFO 的 4 个 tile 预取能掩盖部分 DRAM 延迟,但掩盖不了带宽缺口本身——FIFO 很快被抽干,MXU 只能停下来等。LSTM 更极端:强度只有 64–96,权重等待占比高达 58%–62%,实测性能只有 2.8–3.7 TOPS。

CNN0:唯一贴着峰值跑的应用#

CNN0(源自 Inception V2)的强度高达 2888,远超脊点 1350,落在 Roofline 水平区。它的权重加载等待是 0%——因为卷积权重的空间复用让权重消费速度远低于供给速度,权重管线完全够用。阵列 78.2% 的周期在干活,实测 86 TOPS,达到峰值的 93%。CNN0 证明了 MXU 这套设计本身没有毛病:只要喂得饱权重,阵列是能贴着 92 TOPS 跑的。瓶颈不在阵列,在供给端。

CNN1(AlphaGo):一个值得单独拆的案例#

CNN1 是 DeepMind 的 AlphaGo 网络,强度 1750,理论上应该也是 compute-bound,但实测只有 14.1 TOPS、活跃周期 46.2%、有效 MAC 只有 22.5%(即 65,536 个 MAC 里平均只有约 1.5 万个在干有用的活)。论文归因了三个原因:

  1. 4 个全连接层的强度只有 32:AlphaGo 除了卷积还有 4 个 FC 层,这些层没有空间复用、只有 batch 复用,强度退化为批大小 32,远低于脊点 1350——这几层是硬生生的 memory-bound。论文指出 35% 的周期在等这 4 层的权重
  2. 浅特征深度造成”二维内部碎片”:部分卷积层的输入通道数很小(比如第一层只有 3 个 RGB 通道),展开后填充不满 256×256 的阵列,平均只有一半的 MAC 装着有效权重(有效 MAC 22.5% = 活跃周期 46.2% × 平均约一半的有效载荷)。这正是”稀疏”最原始的形态——不是权重稀疏,而是形状不匹配的浪费,TPU 因为时间紧砍掉了稀疏支持,对这类浪费毫无办法。
  3. 23% 的 RAW 停顿:层与层之间存在数据依赖(下一层的输入要等上一层算完),脉动阵列 512 拍的流水填充延迟在这种依赖链上无法隐藏,转变成显式的等待。

CNN1 的意义在于:它把”memory-bound”和”compute-bound”的二分法打破了——一个应用可以同时被权重带宽(FC 层)、形状碎片(浅卷积)和依赖延迟(层间 RAW)三座大山压着。这也解释了为什么论文说”我们的计数器无法精确解释阵列空闲的每一拍”:重叠执行的流水线上,多种停顿来源混在一起,只能看到部分账目。

负载构成:平均 28% 是被谁拉低的#

最后把账算圆:六个模型里,MLP0 占部署 TPU 的 61%、LSTM0 占 29%(论文 Table 1),两者加起来 90% 的芯片在跑 memory-bound 负载,利用率 8%–13%;CNN 合计只占 5% 的部署量。所以”28% 的平均利用率”根本不是 MXU 设计的失败,而是负载结构决定的——90% 的 TPU 在跑权重带宽喂不饱的模型。这正是论文用 Roofline + 计数器把每一分性能损失定位到具体原因的原因:不是”芯片不够快”,而是”权重带宽不够宽”。

设计空间探索:如果 Google 多给 TPU 六个月#

论文最后做了一个极具价值的”反事实”实验:用性能模型模拟”如果 TPU 有更多设计时间,改哪个参数最划算”。结论在 Figure 11(下图)里,五条曲线分别把 5 个参数从 0.25× 缩放到 4×,看加权平均性能的变化:

论文 Figure 11:加权平均性能对五个设计参数的敏感性。横轴为参数缩放倍数(0.25x 到 4x),纵轴为相对性能(图片来源:TPU 论文 Figure 11)
论文 Figure 11:加权平均性能对五个设计参数的敏感性。横轴为参数缩放倍数(0.25x 到 4x),纵轴为相对性能(图片来源:TPU 论文 Figure 11)

三个结论:

第一,带宽是最值钱的杠杆。内存带宽翻 4 倍 → 平均性能翻 3 倍。细分到负载:MLP/LSTM 从 4× 带宽中获益 3×,但从更高时钟频率中一无所获;CNN 反之,4× 时钟带来约 2× 提升,加带宽几乎没用。这和 Roofline 的预测完全一致——优化方向必须对着你的负载类型

第二,把 MXU 从 256×256 扩到 512×512 反而让性能变差。原因就是上一节 CNN1 里提到的”二维内部碎片”,论文给出了一个精妙的例子:LSTM1 里有一个 600×600 的权重矩阵,

  • 用 256×256 阵列切分:600/2562=9\lceil 600/256 \rceil^2 = 9 步,每步约 2 µs(受 DDR3 加载权重限制),共 18 µs
  • 用 512×512 阵列切分:600/5122=4\lceil 600/512 \rceil^2 = 4 步,但每步要加载 4 倍大的权重块,每步约 8 µs,共 32 µs

更大的阵列让切分更整,但每一步的权重加载时间随阵列面积线性增长,而 DDR3 带宽没变。这就像操作系统的内存分页:页越大,页内碎片越少,但换页成本越高——论文原话是”这个问题类似于大页的内部碎片,只不过它是二维的,更严重”。阵列尺寸不是越大越好,它必须和权重组装线的带宽匹配

第三,TPU’(假想的”多给六个月”版本):论文构造了一个假设的改进版 TPU’——时钟频率提升 50%(700 → 1050 MHz)+ 权重内存换成 GDDR5(带宽提升 5 倍以上,脊点从 1350 左移到 250)。结果:

  • 只加时钟:几乎没变化(MLP/LSTM 不吃时钟);
  • 只换 GDDR5:几何平均性能提升 2.6 倍,加权平均 3.9 倍
  • 两个都做:几何平均 2.9 倍,加权平均不变(还是 3.9)——因为带宽瓶颈解掉之后,下一个瓶颈是主机端。

把主机服务器开销算进去,3.9 倍会缩水到 3.2 倍。代价方面:GDDR5 需要把内存通道从 2 个翻到 4 个,die 面积增加约 10%;但带宽上来后 UB 压力减小,可以把 24 MiB UB 砍到 14 MiB(论文 Table 8 实测六个模型最大只用 13.9 MiB),把面积找补回来;整机功耗预算从 861 W 升到约 900 W。

这个假想实验直接预言了 TPU v2 的设计方向——TPU v2/v3 果然换上了带宽远超 GDDR5 的 HBM 高带宽内存,并且把算力单位从”整颗芯片一个大阵列”改成”多核 + 可编程的向量/矩阵单元”以缓解形状碎片问题。它给整个行业留下的方法论是:推理加速器的性能瓶颈通常不在计算阵列,而在喂权重的那条管子;先画 Roofline 找到脊点,再决定改哪里。这条教训后来被反复验证:现代 GPU(H100 的 HBM3、B200 的 HBM3e)把大半芯片面积和封装预算花在堆内存带宽上,而不是堆计算单元——“带宽优先”已经成了推理硬件设计的第一原则。

局限与未解决的问题#

TPU v1 的脉动阵列 + weight-stationary 组合在 2015 年的推理负载上证明了自己,但它暴露的问题在今天的 LLM 时代依然没有完全解决:

  1. 固定形状的碎片化:256×256 的阵列尺寸是写死的,任何不匹配的层形状(浅通道、非整数倍维度)都会造成 MAC 空转(CNN1 的 22.5% 有效 MAC 就是例子)。今天的 GPU 用更小的矩阵单元(如 NVIDIA Tensor Core 的 16×8×16 MMA 指令)加灵活的调度来缓解,但仍存在类似的填充问题。
  2. weight-stationary 对”动态权重”无能为力:Transformer 的注意力头里,Q、K、V 矩阵是推理时动态生成的,不是静态权重——它们无法预加载驻留。所以现代 GPU 的矩阵单元普遍采用输出驻留/数据驻留混合模式(如 Hopper 的 wgmma 指令支持输出驻留语义),而 weight-stationary 的阵列为每层权重只从内存读一次的优势,在权重动辄几十上百 GB 的大模型上也无法直接延续——权重根本驻留不下,只能分块流式搬运,又回到了带宽问题。把这条线索延伸到现在:LLM 推理的 decode 阶段是逐 token 生成的,每一步都要把所有权重从 HBM 读一遍,还要读整个 KV cache(本站的 MLA 拆解PagedAttention 拆解 里那些压缩与分页技术,本质上都在优化这条读取路径)——权重驻留的思路彻底失效,现代推理优化的主战场恰恰就在 TPU v1 当年卡住的”喂数据”这一端
  3. 编译器负担前置:脉动阵列的确定性执行把调度责任完全推给了编译器(TPU 的用户态驱动/编译器负责决定指令顺序、双缓冲节奏),编译器必须精确知道每个操作的时序。这套”编译期决定一切”的哲学后来被 Groq 的 TSP(Tensor Streaming Processor)发扬光大——Groq 明确宣称”消除所有反应式硬件机制(仲裁器、缓存)“,用确定性流水线换取可预测延迟,并在 ISCA 2020 论文中报告 ResNet50 单图推理延迟低于 49 µs。
  4. 稀疏依旧缺席:CNN1 的”伪稀疏”浪费(形状碎片)和真稀疏(权重/激活中大量为零)是两回事,TPU v1 两者都没解决。到今天,稀疏计算依然是各家加速器(NVIDIA 2:4 结构化稀疏、Cerebras 等)的攻坚方向。

小结#

把三篇核心问题(92 TOPS 怎么来、28% 去哪了、Roofline 怎么用)浓缩成答案:

  1. 92 TOPS 来自 256×256 的脉动阵列:65,536 个 MAC 排成网格,激活从左流入、权重从顶加载、部分和向下汇入累加器,每拍完成 65,536 次乘加;用对角波前 + 双缓冲隐藏内部延迟,让软件看到”一拍读入 256 个输入、瞬间产出 256 个部分和”的幻觉。
  2. 平均 28% 的利用率是负载决定的,不是设计失败:90% 的部署跑在 MLP/LSTM 上,这些模型的操作强度(64–200 MAC/byte)远低于脊点 1350,被 34 GB/s 的 DDR3 权重带宽卡死(权重搬运时间 1.93 µs vs 计算 0.29 µs);CNN0 能跑到 86 TOPS(93% 峰值)证明阵列本身没有短板。性能计数器给出精确账目:平均 43% 的周期在等权重、12% 在移权重。
  3. Roofline 是这一切的定量语言P=min(Ppeak, I×Bmem)P = \min(P_{peak},\ I \times B_{mem}),脊点 = 峰值/带宽。TPU 的脊点在 1350 ops/byte,Haswell 13、K80 9——TPU 用”高而宽”的 Roofline 形状把对手的整个屋顶压在自己脚下,再靠确定性执行把延迟约束下的利用率推到 80%。

这篇讲完,TPU v1 的”为什么设计成这样”和”为什么是这个性能”两座大山都翻完了。下一篇(系列终篇)将回到芯片之外:TPU 的软件栈如何把 TensorFlow 模型编译成 CISC 指令、JIT 缓存机制如何工作,以及”一条指令一层网络”的编译器和调度器设计——那是”15 个月造出芯片”的另一半工程。

参考资料#

  1. In-Datacenter Performance Analysis of a Tensor Processing Unit(ISCA 2017 论文原文,所有数据来源)
  2. H. T. Kung:Why Systolic Architectures?(IEEE Computer 1982,脉动阵列的原始论述)
  3. Roofline: An Insightful Visual Performance Model for Multicore Architectures(Roofline 模型原始论文)
  4. Eyeriss: A Spatial Architecture for Energy-Efficient Dataflow for Convolutional Neural Networks(ISCA 2016,数据流分类与 row-stationary)
  5. Think Fast: A Tensor Streaming Processor (TSP) for Accelerating Deep Learning Workloads(Groq,ISCA 2020)
  6. Google Cloud Blog:An in-depth look at Google’s first Tensor Processing Unit(官方架构解读,含 MXU 脉动阵列图)
  7. NVIDIA:NVIDIA Hopper Architecture In-Depth(现代 Tensor Core 与 wgmma 指令的官方文档)
  8. 博客园:systolic array 架构设计综述(社区解读,三种数据流模式对比图来源)
  9. 腾讯云开发者:【AI系统】谷歌 TPU v1-脉动阵列(社区解读)
  10. Wikipedia:Systolic array(脉动阵列词条)

文章分享

如果这篇文章对你有帮助,欢迎分享给更多人!

TPU v1 完全拆解(二):脉动阵列、weight-stationary 数据流与 Roofline 性能归因
https://pinghaoyang.com.cn/aigc/posts/tpu-v1-part-2/
作者
平昊阳
发布于
2026-08-21
许可协议
CC BY-NC-SA 4.0

评论区

Profile Image of the Author
平昊阳
乘长风,破巨浪, 展鸿图于未央!
--
总访问量
--
访客数
公告
欢迎来到我的个人博客!欢迎关注交流吖!
更多相关公告,见
社交-留言」。
音乐
封面

音乐

暂未播放

0:000:00
暂无歌词
站点统计
文章
66
分类
16
标签
93
总字数
477,284
运行时长
0
最后活动
0 天前

文章目录