课件停在首屏，先不往下滚。

好，我们开始。这一讲叫「TPU 与 GPU」，但我先说清楚， 这门课不给你念参数表。参数表恰恰是我们第一个要拆穿的东西。

⭐ 往下滚一点，停在那道黄色的选择题上。 四个选项先别点 —— 点下去答案就出来了。

不过在开始之前，先回一道上一课的题，热热身。 上一课我们拆的是 DeepSeek V3，总参数 671B，每个 token 只激活 37B。 那 MoE 到底省下了什么？

⭐ 这里一定要真的等。举手、或者让他们喊出来都行， 但要等到有人给出答案再点。 —— 这道题的作用不是考知识，是把房间从「听」切到「答」。 自己念完自己点开，这三十秒就白花了。

答案是「省算力，不省显存」。算的时候只走 37B 那条路，算力是真省了； 可 671B 的权重一个字节都不能少放 ——  你不知道下一个 token 会挑中哪几个专家，所以全都得在显存里待着。

点开答案，指着最后那一段（⭐ 那行）念。

但我真正想让大家听见的不是这个答案，是它背后那件事 ——  显存和算力是两样东西，会各走各的。上一课整整一课都在分开算这两笔账。
而这一课要加的是第三样：带宽。 —— 而且待会儿会看到，这三样里真正决定一个算子快不快的，常常不是前两样。

这句「第三样是带宽」是这一讲和上一讲的接缝， 也是第 1 节的引子。说完就往下走，别在这儿展开 ——  兑换比那个数第 1 节自己会算出来。 ⭐ 再往下滚一点，第二道黄框题。同样，四个选项先别点。

还有一道，这道纯猜。这门课后面会把两颗芯片拆开、一层一层往里看。 在开始之前先猜一下 —— B200 上那个矩阵单元叫 Tensor Core， v7 上那个叫 MXU，干的是同一件事。整颗芯片上，两边各有几个？

⭐ 这道一定会有人猜错，而且错得很整齐 ——  大多数人会觉得两边一个量级。让他们错，这道题的价值全在这儿。

B200 是 592 个，v7 是 4 个。份数差 148 倍。 —— 但真正该惊讶的是下一句：这两颗一拍能做的乘加总数只差 15%。 同样一份算力，一边切成 592 小块，一边切成 4 大块。

⭐ 念完就停，不要解释「为什么」。 指着答案区最后那一行念 —— 数字给你了，「为什么」先欠着。 ⛔ 这里最容易犯的错是当场展开。 一展开，这道题就从钩子变成了剧透， 第 3 节那二十四分钟就没人有问题了。 —— 要的效果是他们揣着「为什么好几百对个位数」听完前两节。 到第 3 节 3.7 那笔账的时候，会有人自己喊出来。

整节课就两个问题，一条路线。第一个问题，这两块硬件在硬件层面 到底哪儿不一样。第二个问题，我们拿同一个算子，从头到尾跑一遍， 看那些不一样分别在哪一步冒出来。那个算子就是 FlashAttention。

往下滚一屏，停在第 0 节那张图。整张图只有两个框， 不用找别的东西 —— 这一讲要做的就是把这两个框说透，然后立刻往下走。

现在看第一张图。全课只有这一个前提，就这两个框。

左边 GPU，出身是图形处理器。它从第一天起要伺候的， 就是一大堆互不相干的小任务 —— 形状不知道，访问模式不知道， 控制流也不知道。所以它的芯片面积，大半花在「应付各种情况」上。 cache 层次、warp 调度器、海量并发线程，全是为了这个。

右边 TPU，出身是为矩阵乘定做的专用芯片。它从第一天就只服务神经网络。 形状是规则的，访问是可预测的，控制流是静态的。 所以那一整套「应付各种情况」的电路，它一样都不用做。

⛔ 这里最容易顺嘴说错，而且我自己说错过： 别说「省下来的面积全给了计算单元」。
省下来的既没变成算力，也没变成片上存储 ——  第 3 节下半会拿两张表把这两条都堵死：乘加总量 GPU 反而多 16%， 片上 SRAM GPU 多 73%。

那它省下来的到底是什么？—— 是复杂度本身。 没有 cache、没有调度器、布局在编译期就定死； 同样多的算力，它只切成四大块，而不是五百九十二小块 ——  少了一百四十八套控制电路。

⚠️ 台下多半会接一句「那是不是更省电」。 这一问有官方答案，而且很硬 —— 但答的时候要带边界。

可以答，念两样就够。第一样是 TPU 初代论文里的原话 ——  它把省掉的东西一口气列完：没有 cache、没有分支预测、没有乱序执行、 没有多线程、没有上下文切换。然后下一句是： 「Minimalism is a virtue of domain-specific processors.」

⭐ 第二样是**一个数**，比什么形容词都管用： 「那颗 die 上，控制逻辑只占百分之二。」
数据缓冲三成七、计算三成、I/O 一成 —— 控制，两个点。 论文紧接着写：这一块在 CPU 和 GPU 上大得多。 ⛔ 然后必须刹车，三条边界一条都不能省：
一，那篇论文比的是 2015 年的 TPU v1 对 K80 ——  那一代省下来确实换成了算力，同样一颗小 die 塞了二十五倍的 MAC。 而到 v7 对 B200，第 3 节下半会数出来 GPU 反而多百分之十六。 同一个机制，十年之间结果反过来了。
二，Ironwood 官方那个每瓦性能翻倍，比的是上一代 TPU，不是 GPU。
三，良率和可靠性 —— 没有出处，不说。 官方讲的是「确定性执行」，那是尾延迟，不是可靠性，别混。 停一拍。下面这两句是这一讲要背下来的。

一个是「不知道你要跑什么，所以处处留一手」； 另一个是「早就知道你要跑什么，所以一手都不留」。

不用滚了，就停在这两个框上说下面这段。

那么这门课要讲的是什么呢。两边处处都不同 ——  但这些不同，全是同一个决定的后果。

哪一个决定？就是把一块数据从 HBM 搬进片上的时候，谁来安排。 GPU 那边有 cache 在运行时替你猜；TPU 那边没有，编译期就排死了。

形状要对齐到 8 乘 128、一条指令吃多大、kernel 谁来写、 连一台机器要付什么代价 —— 全都是从这一处长出来的。

这里语速放慢。这句话是整节课的论点， 后面每一节都在给它找证据。说太快，学员会当成一句口号听过去。

开课前还有一句话我要说死：这门课不比谁强。 两种出身各自都是当年的正确决定，代价也各自都还在。 所以整节课你不会听到「谁更好」这种句子， 只会听到「这件事在这一边由谁决定」。

根立完了。下面直接进第 1 节 —— 它只做一道除法，很快； 做完那道除法，第 2 节就开始拆硬件。

滚到第 1 节。这一节只有两张图，节奏要快。

上一段我说了「只有一个成因」。但在比任何东西之前，得先知道这台机器上「算」和「搬」的兑换比是多少。 这一节就干这一件事 —— 把那个数算出来，不下结论。

两边的官方规格表都给了两个数：每秒能算多少次，和每秒能搬多少字节。 我们把它们除一下。这个比值的意思很实在 ——  搬一个字节要花多少时间，那段时间里这块硬件能算几次。

这句话的方向特别容易讲反，慢一点说。 不是「算 312 次要配 1 个字节的原料」，是「搬 1 个字节的工夫，机器能干 312 次活」。 所以它是一条最低要求：你的算法每搬一个字节，配套得算够 312 次， 才刚好把算力喂饱；算不够，多出来的算力就在那儿空等。 停在那张两栏对照图（fig1-4）。不要急着往下滚， 让数字在屏幕上待一会儿。

左边 TPU v7，一颗芯片 2,307 TFLOPS，带宽 7.37 TB 每秒 （＝ 8 × 921.6 GB/s；本课统一用 7.37，别处写过的 7.38 / 7.4 都是同一个数的取整）。 除出来 312.9。

右边 GB200，NVL72 里那一颗，2,500 TFLOPS 配 8.0 TB 每秒。 除出来 312.5。

这里停死。数字念完不要马上解释， 让台下自己反应过来这两个数是一样的。这是第 1 节唯一的爆点。
⭐ 具体怎么停（2026-09-10 口述后补）：念完「312.5」就闭嘴， 心里默数两拍再开口。「停死」只写了要停、没写停多久， 而临场默认动作是接着往下讲 —— 不写出做法的指令， 到台上就等于没写。（同型教训：外传那条 XLA 报错， 讲义只写「念出来」，可整句英文根本念不出去。）

两家公司，两套架构，两种设计哲学 —— 比值撞在小数点后一位上。 换成 FP8 也一样，625.8 对 625.0。

⛔ 这里千万别说「相对差万分之三」。 那个第三位是假的 —— GB200 那个「8 TB/s」官方没写口径， 十进制和二进制两种读法本身就差 2.4%。
要说的是这一句：「两家、两套架构、两种哲学 ——  算力除以带宽，落在同一个量级上。」 惊讶感一点不减，而且这句经得起追问。
⛔ 这一句是给你自己的判断依据，不要念出来（2026-09-10 口述时 我把它念出去了）。判据：凡是解释「为什么这么讲」的句子都属于讲义， 不属于讲台 —— 念出来会把听众从内容拉到你的方法论上。 指主线上那条蓝线（「同样 8 颗 24 GB 的 HBM3E， 带宽为什么不一样」）。这一句主动讲，别等人问 —— 它每次都会被问。 ⭐ 一句话答完：「HBM3E」只说了代次，没说档位。
每颗 stack 一千零二十四根数据线，带宽就是这个宽度乘以每根引脚的速率。 而这一代从八点零到九点六 Gbps 都叫 HBM3E ——  v7 跑七点二，GB200 跑七点八上下。差别全在引脚速率。
再补一句最要紧的：用哪一档是买方挑的，按功耗、散热、可靠性挑， 不是芯片能力的上限。 ⭐ 先补一句口头过桥，别硬接（2026-09-10 补）： 引出 HBM3E 那问用「说到这儿，有个问题每次都会被问，我主动讲」； 答完转到下面这句用「然后是这一节我最想说的一句话」。
原来这两处是两条舞台指示直接相邻，中间没有可念的句子 ——  台上会出现一个突兀的跳转。
⭐ 然后这一问就把这门课的判据打回到自己头上了。
⭐ 然后把那句最值钱的说出口：「而『这个数是二进制还是十进制』， 正是第 6 节要教的那四句问法之一 —— 它现在打在我自己的招牌数字上。 那四句不是拿来审别人材料的，是先拿来审自己的。」

但接下来这句必须跟着说，不能省：「都落在 312」是真的， 「为什么都落在 312」是我的推断，没有任何一家这么解释过。 而且它挑型号 —— 换成 HGX 板上那颗 B200，就是 281.3，差 11%。

所以我们后面每次拿 312 当判据，用的都是那个带限定的版本。 往外引的时候不能把限定甩掉只留数字。

往下滚一屏，图还没完 —— 下半段有三根长短不一的条。 只念不解释，二十秒。

顺手说一句：刚才那个除法，其实每一层都能做一次。 HBM 是 312，L2 是 119，共享内存是 64。 —— 越靠近计算，这条线越低。

这里不要展开。台下现在还不知道要它干嘛， 解释了也存不住。本节最后讲融合时会回来用它，那时候才有落点。
⚠️ 但要念一句前向指路：「先记着，这一节最后会回来用它。」 （2026-09-10 口述时补的）「只念不解释」如果什么都不交代， 台下会当成一句没头没尾的插话直接丢掉 ——  钩子必须说清在哪儿收，否则它不是钩子，是噪音。 同一条判据在外传那边也踩过：X1 埋的钩子写错了在第几节收。 回到那两个面板 —— 还是同一张图，不用换。

比值一样，还有一种可能：分子分母凑巧成了比例。也不是。 面板上四个数就摆在那儿：算力 2,307 对 2,500，带宽 7.37 对 8.00 —— 两项都在 8% 以内。

这句是这一节真正的收获： 不只是比值一样，分子和分母各自也几乎一样。 凑巧成比例这条退路，堵死了。

顺一句规格表：十行并排看每行都一样重，但值得记的只有一类 —— 一边有、另一边整个没有。TPU 没有硬件自动管的片上缓存那一层， GPU 没有可编程的专用协处理器那一层。缺的不是容量，是「谁做决定」。 —— 那两层具体长什么样，第 3 节有专门的图，这里一个字都不要多说。 完整十行折在本节最后，回查用的手册，现在不展开。

滚到「这条线怎么用」那张三栏图（本节第二张图）， 先停在第 ① 栏。 这一段是本节时间的第二个重点，别当收尾一带而过。

线定出来了，但光有 312 没用 —— 它是硬件那一侧的数， 你还得会算算子那一侧的数，两个数才能比。

算子那一侧的数叫算术强度，算法只有一句话： 数它搬几个字节，数它算几次，除一下。

这句是这一节最实用的一句，说慢一点。 它比 312 本身更值钱 —— 312 会随硬件换代变，这个算法不会。 滚到「这条线怎么用」那张三栏图， 先停在第 ① 栏。算式就印在图上，你只要念着走。

我们当场量两个。

第一个，向量加。y = a + b，bf16。 读 a 两个字节，读 b 两个字节，写 y 两个字节 —— 一共搬 6 个字节。 中间做了几次运算？一次加法。所以强度是 1 除以 6，0.17。

跟 312 比 —— 低了一千八百多倍。

这意味着什么？这行代码在 GB200 上把 8 TB 每秒的带宽全部跑满， 也只有 1.3 TFLOPS —— 那颗芯片峰值是 2,500。 用掉万分之五。算力再翻十倍，它一秒都不会快。

这里可以停一下问：「所以买更贵的卡， 能让这行代码变快吗？」让他们自己答「不能」。 移到第 ② 栏。

第二个，矩阵乘。N 乘 N 的方阵，bf16。 乘加次数是 2 N 的三次方，搬进搬出三个矩阵是 6 N 的平方 字节。 约一下 —— 强度等于 N 除以 3。

那要多大才够到 312？边长接近一千。

反过来说：边长 256 的矩阵乘，强度只有 85，还在带宽那一侧。 —— 「矩阵乘就是算力受限」是个错觉，它得够大才算数。

这一句一般会有人抬头。它推翻的是一个 几乎所有人都默认成立的常识。 同一张图，移到第 ③ 栏。 这四十秒是全课那条线的起点，别赶。

到这儿可以给这门课真正的主角起名字了。

刚才那两行算术其实说的是同一件事：一个算子快不快， 取决于它每搬一个字节做多少次运算。 那想让一个慢算子变快，就只有一条路 —— 把分母做小，少搬。

而少搬最有效的办法，是把连着的几个算子合成一个 kernel， 让中间结果根本不落 HBM。这就是算子融合。

这句停一下：它一个 FLOP 都没省。 分子分毫不动，变的只有分母 —— 很多人以为融合是「少算了」，不是。

所以融合在这条线上的效果，是把算子往右推 —— 从带宽那一侧，推到算力那一侧。

还是第 ③ 栏，指下半那个黄框（312 → 64 那两个数）。 这一段是补上去的，因为上面那句话会被听岔。

刚才我说「中间结果根本不落 HBM」。这句话很容易被听成 「中间结果消失了」。—— 它没有消失。

它被塞进了片上暂存：GPU 上是那块最多 227 KB 的共享内存， TPU 上是 VMEM。HBM 那一层确实少搬了，可那批数据总得在片上落脚。

这句停一下：搬运没有消失，是换了一层。 刚才那三根条，用处就在这儿。

而片上那一层有它自己的线 —— 刚才念过，是 64，不是 312。 所以融合完要重新对一次账：拿新的强度跟 64 比。

这里必须多说三句，否则上一句自相矛盾： 门槛从 312 降到 64，是降不是升。

那为什么还会撞墙？因为动的不只是门槛，还有分母。 融合把 HBM 那一层的流量转成了片上流量 —— 片上搬的字节数能涨几十倍， 而门槛只降了 5 倍。撞墙不是因为门槛变高，是因为分母变大。

台下常会接一句：「那干脆留在寄存器里不就完了？」 —— 大多数时候做不到。一个 SM 的寄存器堆总共 256 KB， 跟共享内存一边大；而且你多占一份，能待命的 warp 就少一批。 这笔账第 3 节那张 SM 拆解图上有。

1.2　方阵只是特例 —— 强度由最小的那个维度说了算（图 1-8）

滚到 图 1-8。⏱ 这一段 3 分钟， §1 相应地从 10 分钟变成 13 分钟。

好，线定出来了，也会用了，也知道它是干什么用的。 第 1 节的全部产出就一样 —— 一个数，和一个算它的方法。 数会随硬件换代变，方法不会。

⛔ 2026-09-05 起课件上那个「产出就两样」的框已经删了 —— 原来第 ② 条讲「一边有、另一边整个没有才值得记」， 那条判据在 §1 是悬空的（本节根本没推出它），而它在 3.5 / 3.6 用到的地方 上下文里自明，不需要预先定义。 台上也别补讲，讲了就是把删掉的东西又搬回来。

—— 接下来每一节都在这条线上。第 2 节看融合前后差多少： 同一个注意力，两种写法 FLOP 一模一样，强度差一千倍； 第 3 节看它在两边分别怎么落地； 第 5 节回答那个真正的问题 —— 这件事该由谁来决定。

滚到第 2 节开头，停在那张访存路径图（fig2-1）。 这张图是全课的中心，值得在屏幕上多停一会儿。

上一节的结论是两边胃口一样大。这一节说的是 ——  吃法完全不同。

⭐ 第一个决策点要单独解释一句，别当术语滑过去。
「coalescing 就是合并访存：一条 load 发出去，warp 里那三十二个线程 各算出一个地址，硬件得把它们合并成尽可能少的内存事务。」
再补一句量级：合得好不好，差八倍带宽 —— 而且只有跑起来才知道， 因为编译器不知道你的地址是怎么算出来的。」 ⛔ 这一讲最容易讲砸的方式：讲成「TPU 没有 cache」。 那是把一整套机制缩成一个零件 ——  而且第一个决策点根本不是 cache，是 warp 那三十二个地址能不能合并。
要交的是「4 对 0」，不是「有没有 cache」。

我们不看任何参数，就跟着一次取数走一遍： 从 HBM 里拿一块数进片上。这一路要做几个决定，每个决定是谁做的、什么时候做的。

先只讲左半边（GPU）。一个决策点一个决策点地指， 每指一个说一句「这个答案，编译的时候不存在」。

左边 GPU。一个 warp 三十二个线程，各给出一个地址。 第一个决策点：这些地址能合并成几条 cache line？—— 跑起来才知道。

接着走到 L1，第二个决策点：命中还是没命中？再走到 L2， 第三个：命中还是没命中？没命中要换出一块，第四个：换谁出去？

四个。而且注意 —— 这四个我一个都没说「慢」。 红色的意思不是慢，是「跑起来才知道」。

切到右半边（TPU）。这边节奏要快，因为它没有东西可数。

右边 TPU。编译期：XLA 决定每个数组切成 8 乘 128 的块。编译期：决定什么时候搬、 搬多大一块。编译期：算准 VMEM 装不装得下。然后运行时只剩两件事 ——  DMA 整块搬过去，向量单元直接取。

路上零个运行时决策点。而且下面那一格是空的：没有 cache 这一层 —— 所以也就没有命中、未命中、替换这些概念。

这里停死，念图上那句红字：「四个，对零个。」 这是全课的中心，第 0 节那个论点在这一刻第一次落地。 念完不要解释，让它挂在屏幕上。

那绿色是什么意思？绿的也不是「快」，是「写代码的时候就已经定死了」。

所以两边的分工从这里就定了。GPU 那一侧，把这几个决定交给硬件在运行时处理 —— 也就是交给了人，交给写 kernel 的那只手，你能做的是让硬件更容易猜对： 访问连续、避开 bank、手工 tiling。
TPU 那一侧，交给编译器在编译期算准 —— 也就是接受它的脾气。

翻回第 0 节那个绿框「这门课不比谁强」，或者直接说这一句： 这不是谁更聪明 —— 是两条路线，各自靠一样东西。 一定要说，它挡掉后面所有「所以谁更好」的提问。

GPU 靠运行时适应 —— 形状对不对都能跑起来； 代价是命中率要跑起来才知道，不好的时候只能反复试。
TPU 靠编译期算准 —— 形状提前定死，一个决策周期都不浪费； 代价是运行时没有补救手段，形状不对只能回去改代码。

⭐ 这里加一句，它是 §0 到 §2 那条线的收口： 「到这儿，『处处留一手』才第一次不是一个形容词 ——  它就是这四个决策点，每一个都是真实存在的电路， 这一次的地址合不合并、命不命中，它都在那儿等着被问。」 —— 开场那句抽象的话，在这一刻变成了硬件成本。
⛔ 说到这儿就停，不要接「所以省下来的面积给了谁」 —— 那个方向本课撤回过，没有出处。

—— 后面第 3 节、第 5 节看到的每一样东西，都是这一句的后果。

⛔ lane / sublane 这一段不在这儿讲了 —— 它搬到了 3.3 开头（那儿才第一次真用上）。 这一讲到这里就收，直接滚到 2.3。 滚到 2.3，停在 Flash 那张对照图（fig2-3）。

判据有了，坐标有了。现在挑一个算子跑到底 —— FlashAttention。 理由很实际：它是今天每个大模型都要跑的那一步， 而且它正好卡在 312 的错误一侧，所以两边都必须动手。

先记一个数。一个头，128K 序列，bf16，那个注意力分数矩阵是 整整 32 GiB。

给一个对照再停一下：一个 v7 device 一共只有 94.74 GiB 可分配 —— 装不下三个头。而模型有 128 个头。

朴素写法是怎么走的？算出分数矩阵 S，写回 HBM；读回来做 softmax， 再写回去；再读出来乘 V。一个头，来回搬了四趟 32 GiB —— 而真正的乘加，只有那两次矩阵乘。

Flash 怎么做的？把 Q、K、V 按行切成块，一次只取一块。 这一块的 S 只在片上出现，不落地；softmax 改成 online 的， 只维护两个标量 —— running max 和 running sum；累加进输出块， S 当场丢掉，换下一块。

结果是：那个 seq 乘 seq 的矩阵，一次都没在 HBM 里出现过。

指第 ③ 栏。这一栏是这一节和第 3 节之间的桥。

第 ③ 栏问的是「那一块该多大」。两边问的是同一个问题，答案交给不同的人。
GPU：kernel 作者自己定，预算是 228 KB 的 shared memory， 块大小是个可调超参 —— 调错了也能跑，只是慢。
TPU：编译器定，预算是 64 MiB 的 VMEM，但形状必须对齐 8 乘 128 —— 对不齐就补零，而补出来的零照样占 VMEM、照样过 MXU。

滚到最后那张 312 的轴（fig2-4）。

现在把这笔账放回上一节那条线上。

朴素：算术强度 64。它每搬 1 个字节只算了 64 次； 而搬这 1 个字节的工夫，硬件本来能算 312 次 ——  差了近五倍，算力有八成在空等，死死卡在带宽上。

Flash：65,536。同样的 FLOP，搬的字节少了一千零二十四倍， 直接跨到右边，卡在算力上 —— 而那才是该卡的地方。

⚠️ 这句要跟着说：65,536 是理想上界。真实的 Flash 会重复读 K 和 V， 到不了这个数。但结论不变 —— 它是从 312 的左边跨到了右边。

先别急着滚过去 —— 这张轴下面新加了一个绿框 「那它过了片上那道门槛吗」，两分钟，讲完再进橙框。

刚才说 Flash「跨到右边」，跨过的是 HBM 那条线 —— 312。
可片上还有一条：共享内存是 64。

⚠️ 这里必须主动点破，不要等人问： 屏幕上现在有两个 64 —— 轴上那个是朴素注意力的强度， 新框里这个是共享内存那一层的门槛。 一个是算子的，一个是硬件的，数字撞上纯属巧合。 不点破，台下一定当成同一个数。

而融合干的事，恰恰是把字节从 HBM 挪到片上。 —— 所以那道除法得拿新的分母，再做一遍。

指绿框里那五行推导，一行一行念，不要跳。

设块的行数 Br、列数 Bc，head_dim 是 d。
每走一步，从共享内存读进 K 块和 V 块，各 Bc 乘 d。Q 那一块一直待在寄存器里不动。
算的是两次矩阵乘，4 乘 Br 乘 Bc 乘 d 次。搬的是 4 乘 Bc 乘 d 字节。

约一下这一步停一秒再说答案： 片上强度就等于 Br —— 块的行数。 让它自己浮出来，比直接报数有用得多。

Br 取 128，就是 128 对 64，过关。
Br 取 32，就是 32 对 64 —— 在片上重新变成带宽受限。

⭐ 然后把锚点给他们：真实的 FlashAttention kernel， 块行数一律取 64 或者 128。不是巧合，是这道除法逼出来的。

⚠️ 限定要说：这笔账只数了共享内存那一层， softmax 的向量运算和寄存器那一层都没算。它给的是量级，不是精确值。

⚠️ TPU 那一侧的分子是 1,153.5，不是 2,307 —— VMEM 是每个 core 一块， 分母是 per core，分子就得跟着取 per core。这就是那四句问法的第三句， 说错一倍。

最后那行小字（TPU 侧）照念，一个字都不要加。

TPU 那一侧算法完全一样，分子取 1,153.5（per core），分母换成 VMEM 带宽 ——  而这个数官方没有公开，我们不推测。
能说的只有方向：片上带宽比 HBM 高一个量级，门槛就低一个量级， 落到几十这一档，和 GPU 那个 64 同一个数量级。
—— 再加上 TPU 的工作台是 64 MiB，卡住它的不是这条线，是 3.4 那个刀宽。

滚到本节最后那个橙框「回头看一眼」。

好，回头看一眼刚才发生了什么。Flash 做的事， 就是把「算分数 → softmax → 再乘 V」这一串合成了一个 kernel —— 这就是上一节起的那个名字：算子融合。

这句停一下：它连一次乘加都没有省。 图上第 ①② 栏是一模一样的。省的全在分母。

所以到这儿，故事的前提齐了：判据 312，朴素是 64， 所以必须改写。

而真正的问题不是「要不要融合」 ——  是「这个融合谁来做」。Flash 这一次，是人写出来的， 一篇论文加一堆手写 kernel。第 5 节要问的就是：非得靠人吗。

滚到第 3 节开头，停在 GPU 整颗那张（ms-g1）。 这一节十四张图，一开始就要立好节奏 —— 前面走快，钱花在 3.6。

前两节我们说了「两边胃口一样大」，也说了「吃法完全不同」。 这一节把两颗芯片真的拆开，看那个不同长在哪儿。

先看整颗。B200 一个封装里是两颗 die，中间用一条十 TB 每秒的桥连起来。 但你在 CUDA 里看到的是一个设备 —— 缝被藏起来了。

切到 TPU 那张（ms-t1）。两张图要连着指，它们是一对。

TPU v7 也是两颗 die。但它如实暴露成两个 device，两半各有自己的 九十六 G HBM，地址空间互不可见 —— 想用对面的，你得自己写一步。

停一拍：「缝在明处，还是缝在暗处。」 这就是第 2 节那个决定在封装这一层的样子。

藏起来的那一边，代价不会因为藏起来就消失 ——  你付了钱，只是不知道自己付了。跨 die 那一趟，带宽从二十一掉到十六点八， 数字在 3.2c 那张表里。

⚠️ 图上那个 131,072 先别急着往峰值上乘。
它数的是格子 —— 两个 MXU 一共两乘二五六平方个。
而每个格子一拍做两次乘加，这一步图上没有。

⚠️ 这句必须主动点破，别等人问： 「两次乘加」不是「一个乘加一个加、一共两次」。
一次乘加本身就是一乘一加 —— a 乘 b 再加 c，一个周期里一起做完，算两个 FLOP。
所以一格一拍是四个 FLOP，不是两个。 这一步听岔了，后面整条链就差一倍。 在白板上写：1 次乘加 ＝ 1 乘 ＋ 1 加 ＝ 2 FLOP；2 次乘加 ＝ 4 FLOP。

把错算式念出来，别只说「会算错」： 131,072 乘 2 个 core 乘 2 FLOP 乘 2.2 GHz，得 1,153 ——  刚好是图头那个 2,307 的一半。少的就是「每格两次」那一步。

⭐ 如果有人当场追问「那正确的链呢」， 3.4 现在补了一句桥，可以先给他一句预告： 「图上那个是格子数，3.4 那个 262,144 是每 MXU 每周期的 FLOP ——  一除就是每格 4 FLOP，也就是两次乘加。」 不要在 3.1 就展开算，节奏会塌。

往下滚，SM 那张（ms-g2）。

再往里一层。先看 GPU 有什么，再看 TPU 没有什么 ——  顺序不能反，反了就会以为 TPU 是个简化版。

一个 SM 里：warp 调度器、二百五十六 KB 寄存器堆、 二百二十八 KB 的 L1 加共享内存、记分板。 这些部件加起来只干一件事 —— 在运行时动态地藏住延迟和分支。

切到 TPU 那张（ms-t2）。指右边那五个空格。

TPU 这边，这五样一样都没有。

但重点不是「少了五样」，是这五样其实是一件事。 要不要在运行时换任务跑？答「要」，你就必须同时驻留几十份上下文、 必须有东西挑下一个、必须知道谁的数据到了 —— 五样缺一不可。 答「不要」，五样同时失去存在的理由。

这句是这一小节的落点，值得停： 「有没有 cache」和「有没有 warp 调度器」不是两件事， 是同一个决定的两个面。

—— 而第 2 节主路那四个红框，就是答「要」之后的必然产物。

滚到那张十行 SRAM 表。这一段没有图，全靠表。 语速可以快，但三个口径那个框要停。

这里必须正面接一个反问：「你说 TPU 把面积给了暂存， 可 GPU 片上 SRAM 明明更多啊。」

—— 他是对的。我们把两边的片上 SRAM 一格不漏地列出来： GPU 二百三十一 MiB，TPU 一百三十四 MiB，GPU 多百分之七十三。 这一条别嘴硬。

承认完，再翻面。这个顺序很重要 ——  先认下来，后面那两个口径才听得进去。

但同一张表可以读出三个方向不同的结论，看你在问什么。

第一个口径：一共多少 SRAM —— GPU 赢。
第二个：只数「软件当数据暂存用的那一行」——  TPU 一百三十对三十七，约三点五倍。
第三个，也是真正有用的那个：一个计算单元能当成一整块连续工作台用的，有多大？ —— TPU 六十四 MiB，GPU 二百二十七 KB，二百八十九倍。

指那句「GPU 那 37 MiB 不是一整块」。

为什么差这么多？因为 GPU 那三十七 MiB 不是一整块， 是一百四十八个互相看不见的小岛，每个二百五十六 KB， 而里面 L1 还占着一部分 —— 一个线程块能显式管到的最多二百二十七 KB。

这里停死。「二百八十九倍」这个数， 3.6 决定两边把块切多大的时候会直接用上。

还有一句必须说：第一个口径虽然对，却最容易把人带沟里。 GPU 领先的那九十七 MiB，几乎全来自 L2 那一行 —— 而 L2 是这张表里唯一一格你没法安排的空间。 你不能把一块权重「放进 L2 待着」，只能发访存指令，然后希望它还在。

滚到带宽那张表。这一段只讲一个反直觉，两分钟。

上面回答的是「有多大、归谁管」，还差一半：跑多快。 补上之后会撞出一个几乎所有人都猜错的结果。

大多数人默认「cache 总该快一个数量级吧」。B200 的 L2 不是。 HBM 八 TB 每秒，L2 只有二十一，两点六倍。

停一下，让这个数落地。二点六倍。 ⛔ 这里马上要补一句，不然「只有二点六倍」会被听反。 台下很容易得出「片上存储没什么了不起」—— 完全相反。

二点六倍答的是一个很窄的问题：这一个字节，命中 L2 比去 HBM 快多少。
换个问法就完全不同：每一字节容量能供多少带宽？ —— L2 是用一千五百分之一的容量做到二点六倍带宽的， 两者一乘，它的带宽密度是 HBM 的四千倍。

⭐ 给一个能想象的说法，比倍数好记： 「把 HBM 整个读一遍要二十四毫秒；把 L2 整个读一遍，六微秒。」 ⭐⭐ 然后是这一段真正要留下的一句： 「这两个数的比值，恰好就是容量比 ——  也就是说，片上存储的本事全在『它小』这件事上。 它快不是因为它是 SRAM，是因为它小、离得近、线短。」
⛔ 说完立刻拉回来，别让它盖过主线： 「但优化的时候还是只能用二点六倍那个数 ——  因为你搬的是字节，不是容量。」

所以 L2 的价值不在「它本身快」，在「省掉一趟 HBM 往返」。 命中了，你拿到二点六倍；没命中，你什么都没省，还多绕了一圈。 —— 这就是为什么「让它落在 L2 里就好了」不是一个能依赖的策略。

为什么慢？因为 L2 是全片共享的，数据得穿过片上互连才到得了 SM。 而共享内存四点三倍、VMEM 高约一个数量级 ——  它们就贴在计算单元旁边。

这句是这两小节合起来的落点： 「真正快的那两块，恰好都是软件说了算的那一类。」

⚠️ TPU 那一行我们只给量级，不给数 ——  VMEM 带宽是未公开规格，官方只给容量。 能说的到此为止：它比 HBM 高约一个数量级。 别往这儿填一个精确数，也别说成某个整数倍 ——  HBM 带宽是公开的，写出倍数等于把那个未公开的数说出来了。

指最后那个 ⛔ 框「别把这两小节读成 TPU 赢了三局」。 这个框一定要念，它是整节课的护栏。

那三条领先都是真的，但每一条都有对价。
① 暂存又大又快，代价是没有兜底 ——  GPU 猜错只是慢一点，TPU 编译器切错了就是真的停在那儿等。
② GPU 那两个三十七 MiB 看着「没用在数据上」， 它买的是别的东西：形状不规则、长度会变、分支不确定的活，照样能跑得不太难看。

—— 所以真正的结论不是「谁的片上存储更好」，是这一句： 两边都在处理同一个东西 —— 不确定性，只是把它放在了不同的地方。 GPU 花硅去兜住它，TPU 把它消灭在编译期。

滚到延迟阶梯那张图。三到四分钟。这是上半节的落点，别赶。

容量讲完了，带宽也讲完了。还差第三个轴：等多久。 而这个轴，是三个里面唯一一个 GPU 自己也答不上来的。

⭐ 先把问题原样抛给台下，别急着给数： 「一个 warp 要一块数据，L1 里没有，去 L2 查，L2 也没有，再去显存搬 —— 这一串下来要多久？」

看图。L1 命中，三四十拍。L2 命中，二百七十三拍。一路到显存，六百五十九拍。 —— 同一条 load 指令，在源码里长得一模一样。

停一下。三十拍到六百五十九拍，二十倍。

⛔ 这里要拦一个几乎人人都会犯的算法： 「L1 查一次，加 L2 查一次，加 HBM 搬一次」—— 不能这么加。 六百五十九是一次全没命中的端到端耗时，前面那两次扑空已经含在里面了。

⭐ 然后是这一小节真正要问的那句： 「每种情况耗时都不一样，那这东西怎么提前规划？」
—— 答案很干脆：规划不了，所以 GPU 根本不规划。

它的三条对策，没有一条在试图算准。
第一，超量线程 —— 一个 SM 上驻着几十个 warp，这个卡住就切下一个。 延迟不是被消除的，是被别人的活盖住的。
第二，记分板 —— 硬件在运行时盯着操作数到没到，到了才发。
第三，缓存本身就是一台猜测机器 —— 它赌你还会再用一次。

TPU 走的是反方向：把不确定性从源头删掉。 没有硬件缓存，就没有缓存缺失，也就没有可变延迟。 VMEM 里的数在那儿，不是因为运气好，是因为编译器自己发的那条 DMA 把它放在那儿的 —— 而且它也知道该在第几拍等。

⭐⭐ 这一句是上半节的落点，念慢一点： 「GPU 用『总有别的活可干』来盖住不确定的延迟； TPU 用『把不确定性删掉』来避免它。 一个是统计学的答案，一个是确定性的答案。」
—— ⭐ 紧接一句：「所以 TPU 没有 L2 不是缺功能。 缓存是一台猜测机器，而一台编译期就把每一拍排好的机器，不需要猜。」

⭐ 顺手送台下一条实用的： 这也解释了为什么两边的「调优」没有一样东西是重合的。 GPU 上你调的是「怎么让别的活足够多」—— 占用率、warp 数、访存合并； TPU 上你调的是「怎么让编译器排得开」—— 对齐、分块、双缓冲。 同一个词叫调优，动的旋钮一样都不挨着。

⚠️ 最后别让它变成「两边越走越远」： 「现在追峰值的 CUDA kernel，异步拷贝、显式 barrier、自己管双缓冲 —— 那就是在 GPU 上手写一个静态调度。 一旦你真的要峰值，猜就不够用了。」
⭐ 这句给 3.5 那几个框埋了扣，讲到那儿会回收。

⚠️ 数的口径要交代一句：这组数是 H100 和消费级 Blackwell 上量的， B200 没有同类公开数据。量级和形状可以借，别当成 B200 的实测值。

先停在 3.3 开头那张寄存器形状图（fig2-2）。 一分钟讲完，它是词汇表不是内容 —— 但后面 3.3 和 3.4 全靠它说话。

先插一对词。下面要谈「一个线程其实是一条什么」，还有 3.4 那个「八乘一百二十八」， 都得先有 TPU 那边的坐标词，而这两个词的字面会骗人。

两步就能把画面立起来。
第一步，先只看一条：一个 sublane 就是横着的一条， 128 格那么长，512 字节 —— 这是 VMEM 的架构字长，不是凑出来的数。
第二步，再叠 8 条：就是一个向量寄存器，8 × 128 ＝ 1,024 格。

两个词到这儿各归各位：lane 量的是长度，第几格，0 到 127； sublane 量的是第几条，0 到 7。
—— 「对齐 8 乘 128」就是大白话：竖着凑够 8 条，横着凑够 128 格。

⚠️ 名字骗人的地方在这里：sublane 不是 lane 的一小段， 它是横跨全部 128 条 lane 的一整条。「sub」说的是第二维，不是更小。

这里可以轻松一点：「几乎所有人第一次都理解反， 包括我。」承认自己也搞反过，比强调「别搞反」有用。

记法只有一句：最内维是 lane，次内维是 sublane。

然后是这张图真正要留的东西：XLA 故意把内存布局做成跟寄存器一样的形状。 搬进来就能直接算，一次重排都不用。

⛔ 这里有个说法我讲错过，别照老版本念： 「一行在内存里不再连续」—— 那不是代价，那正是目的。 一段一百二十八个元素正好铺满一百二十八条 lane，这是要的效果。」

真正的代价是另一样：这个形状被钉死了。 三笔，都能点名 —— 换最内维要真搬数据（分块绑在最后两维上， 转置和某些 reshape 不是改一下步长，是物理搬运）；不整除就补零， 补出来的零占 VMEM、过 MXU、还算进 FLOP；跟外部交接要重排。 这笔钱第 3 节要付。

⭐ 顺手接一句对照，它把这一段接回主线： 「GPU 那边的合并访存也想要连续 —— 但那是软约束， 没对齐照样跑，只是慢。这边的分块是硬约束 ——  形状不对，要么补零，要么插一个重排。一边罚款，一边直接不让过。」

这一节先不要滚图。头两分钟只说话， 推到「六层」再上第一张图。

刚才少掉的是硅片上的部件。往上一层，编程模型里也会少掉东西。 同一个原因，第二次露面。

这一节我先归到一个问题上。

停一拍，这句要慢：「你写 GPU 代码的时候， 不知道这堆活会落在哪个 SM 上，也不知道什么时候落 ——  是硬件在运行的时候才决定的。」 白板上就写四个字：「运行时才定」。

既然不知道，你就没法点名。写不出「让三号 SM 干这个」。

你能做的只剩一件事：描述关系。 哪些活必须待在一块儿，哪些落哪儿都行。

「必须待在一块儿」的那一批，就是一个 block。

为什么是「必须」？因为这批活之间要互相传中间结果。 而中间结果放在 SM 自己身上那块小内存里，不在外面。
所以一批活一旦落到某个 SM 上，就不许再搬走 —— 搬走了，中间结果就找不着了。

反过来，互相不用说话的活，落哪儿都行。全部加起来，叫 grid。

这两层是你说了算的。 你写代码，其实就是在说两句话：这些必须在一起，总共这么多。

硬件那边还有它自己的分组，理由完全不一样，是为了省电路。

取一条指令很贵。那就让三十二份数据共用同一条指令，一起算。 这三十二个绑成一捆，就是 warp。这个数你改不了，硬件焊死的。

到这儿四层齐了：底下两层硬件定死，上面两层你定。

剩下两层是后来补的，各被一个很具体的麻烦逼出来。
warp group —— 新一代的矩阵指令，要四捆一起发才凑得出形状。
cluster —— 相邻几个 SM 想直接够到对方那块小内存，不想绕远路。

⭐ 这里补一刀，效果很好： 「而『绕远路』是绕哪儿？就是绕那个只比显存快 2.6 倍的 L2。」 —— 3.2c 那个反直觉的数，在这儿变成了一整层编程模型。 前面埋的数在后面变成结构，台下是听得出来的。 现在才滚到第一张图（每一层被什么逼出来）。 指「被什么逼出来」那一列，别念规模那一列。 这一节的落点，说完停两秒： 「这六层，全是从刚才那句『不知道』长出来的。」 ⛔ 图注最后一句已经是「＝运行时适应的组织成本」， 别再念一遍，念了就是同一句说两次。 滚到 ms-t3。这一段指的是空格，不是格子里的字。

而 TPU 那边，编译的时候就已经把一切定死了，不存在「不知道」。 所以这一整套语言，它一层都不需要。

⚠️ 顺手把一个陷阱按住：两边的层根本不在回答同一个问题。 GPU 那些层是用来描述「谁必须在一起」的，TPU 压根不需要描述这件事。 所以这张按问题排，不按名词排。

空的第一格：没有「一组线程」。GPU 有 thread block；TPU 没有这个中间概念。
空的第二格：没有「换一个跑」。GPU 有四个 warp 调度器、六十四个 warp 槽； TPU 一个都没有。

—— 空的那两行，正好是 3.2 里少掉的那五样在编程模型上的投影。 硅片上少了部件，代码里就少了概念。

⭐ 三处「不对齐」里，现场只讲第一条，另外两条留给提问： 「一个 SM 上同时住着好几个 block，它们挨着，但互相看不见对方那块小内存。 第 5 节讲的『占用率』，量的就是这个 SM 还能塞下几批。」 滚到「一个 SM 拆开」那张四格图。

再往里拆一层。「一个 SM」到这儿为止一直是个黑盒。 里面是四个格子，每个格子自己一套调度器和寄存器。 三十二是一个格子的宽度；一百二十八是四个格子各出一捆。

这里有个问题值得单独说：一个线程到底独占什么。

算的部件，它一样都不占。 CUDA core 也好、Tensor Core 也好，都不是分给谁的 ——  它们是流水线上的一个位置，哪一捆这一拍发了指令就用这一拍，下一拍换别人。

它真正独占的只有寄存器。一批活落到 SM 上就划走，跑完才还，中途从不搬动。

⭐ 这一条同时解释了两件事，两句连着说： 「正因为寄存器不搬动，换一捆上来才是零成本的； 也正因为独占，先卡住你的一定是寄存器不够，不是算的部件不够。」

⛔ 三十二位、四个字节、上限二百五十五、按八取整 ——  这些全在折叠里，现场一个都别念。

滚到最后一张（三个都叫「线程」）。

最后回到最底下那一格。一捆是三十二个线程， 每个线程在这一捆里有个编号，零到三十一。那个位置，就是一条 lane。

所以 GPU 上也有 lane，硬件上就是一条 SIMD 通道 ——  跟 TPU 那边的 lane 是同一样东西。

⭐ 这一节最后一句，说完就走： 「差别只在一件事 —— GPU 的 lane 有名字，TPU 的没有。 GPU 那条你能对它单独写 if；TPU 那条你写不出『第三十七号 lane 干点别的』。 所以 CUDA 说的线程，是一条有名字的 SIMD lane。」

还停在 3.3 那张「线程其实是一条 lane」的图上，别往下滚。 ⛔ 先主动打掉一个画面，台下十个有九个脑子里是它： 「一个线程负责输出矩阵里的一格。」
这个画面在没有 Tensor Core 的年代是对的，在今天是错的。

先讲对的那个年代。一个线程就是输出矩阵里一格的所有者 ——  它自己跑那个 K 循环，从共享内存读 A 的一段、B 的一段进自己的寄存器， 乘加，累加在自己手上，最后写回。搬运和计算，它都干。

⭐ 然后把它推翻，这一句要慢：
「用上 Tensor Core 之后，矩阵乘指令是整个 warp 一起发的一条指令。 它的操作数散在三十二个线程的寄存器里，每人拿一小撮。 单个线程，自己算不出任何一个完整的输出元素。」
要具体就给一条：一条典型的指令里，每个线程交出 A 的八个数、 B 的四个数、累加器四个数。就这些。 ⭐⭐ 这一段的落点： 「所以『一个线程负责哪一格』这个问题，到 Tensor Core 上就失去意义了。 线程退化成两样东西 —— 一份寄存器的持有者， 和发那条指令的三十二分之一。」
—— 这才是「一个线程其实是一条 lane」的完整含义。 指下面那个「到底是不是 SIMD」的框。

这一问答案是分层的，两个层面结论相反。 硬件上是 SIMD —— 一个 warp 的三十二条 lane 共享取指、译码、发射， 同一拍执行同一条指令。编程模型上不是 —— 你写标量代码， 而且 Volta 起每个线程真有自己的程序计数器，可以各走各的分支。

⭐ 一句话收： 「SIMT 就是：用标量的写法，跑在 SIMD 的硬件上。 代价写在分支上 —— 真发散了，各条分支只能轮流走。」
然后接回主线：「而这一整套，就是『靠运行时适应』要付的钱。 TPU 不做这个决定，所以它那边连『线程』这个词都没有。」 ⚠️ 被问到搬运再展开，别主动讲（会吃掉两分钟）： 线程管不管搬运？管，但这一半正在退场 ——  Ampere 之前线程自己搬、搬的时候被占住；Ampere 起 cp.async 绕过寄存器； Hopper 起 TMA 一个线程发一条描述符就搬一整块，另外三十一个不参与。

滚到 3.3b，先指主线上那条蓝线的一句话定义 （「记分板就是一张表……」）。这一句必须念，十五秒。
底下那个折叠是机制（四分钟，全讲最容易超时的一段）——  台下没人追问就别展开，直接跳到 3.4。 这一句念出口就行，别加料： 「记分板就是一张表，记着每条指令要的操作数到齐了没有； 调度器每一拍扫一遍，谁齐了让谁发射。」
—— 然后立刻接回主线：「TPU 那边这张表整个不存在。」

刚才说了两遍「GPU 把 lane 暴露给你，代价是一整套东西」。 这一小节就回答一句话：那一整套，具体是哪几样。 三样，各管一个很窄的问题。

三个问题先一口气报出来，报完再逐个展开 ——  学员知道一共几件、现在在第几件，才不会中途掉线：
「一，谁现在能发指令；二，谁跟谁怎么对齐； 三，数据怎么在不占用线程的前提下就位。」 最后一件事，指着 3.3b 那张三代对照图。这是这一小节的钩子，别省。

线程管搬运的那一半，正在退场 —— 三代，一代剥一层。

Ampere 之前，线程自己搬：从全局到寄存器，再到共享内存。 搬的时候它就被占住了，什么也干不了。
Ampere 起，多了一条指令，让数据绕过寄存器直通共享内存 —— 线程不用当二传手了。
Hopper 起更狠：一个线程发一条描述符，就能搬一整块， 另外三十一个完全不参与。

⭐ 这里要把三代收成一句，慢点说： 「三代下来，线程从『搬运工加计算工』，被剥成了一个寄存器的容器。」 —— 它跟前面那句「线程算不出任何一个完整的输出元素」是同一件事的两半： 一半剥走了计算，一半剥走了搬运。 ⛔ 不要在这儿接「所以 GPU 在往 TPU 靠」。 这条趋势线是 §5.4b 的，那里有完整的证据和一条重要的限定 （挪的是「怎么搬」，不是「谁决定」）。这儿只留钩子。

—— 而线程被剥空之后，真正决定「一次能吃多大一口」的就不再是线程了， 是指令。那一维两边差 16 倍，3.4 会当场量给你看。

但先别跳过去。下一小节 3.3b 要先把账单打开。

上面这一整套 —— 让你写标量、硬件替你凑成 SIMD、发散了替你排队、 再替你在几十批活之间来回换 —— 「靠运行时适应」这五个字，到底要花多少钱？ 下一小节把它拆成三套具体的电路，一样一样看。

第一样，记分板。先送一个记得住的钩子： 这个名字是从一九六四年那台 CDC 6600 借来的 ——  那是第一台指令能不按顺序完成的机器，它里面那个部件的正式名字就叫 Scoreboard。 （为什么挑这个字、以及它和现在这块已经不是一回事， 折在课件里，别现场展开。）

回到 GPU。调度器每一拍要挑一个「就绪」的 warp 发指令 ——  「就绪」是谁判的？NVIDIA 把指令按一条线劈成两半： 编译器知不知道它要跑多久。
知道的那一半根本不用硬件管，编译器在指令里直接写「停几拍」。
不知道的那一半 —— 访存、特殊函数、矩阵指令 —— 才是记分板： 每个 warp 六个计数器，发出去加一、写回来减一； 等它的那条指令带个掩码，写明要等哪几号归零。

停：「所以它不是一张大表，是六个很小的计数器。 它只回答一个问题 —— 这个 warp 现在能不能发指令。」

第二样一句话带过：SM 里一小组硬件 barrier，warp 陆续到达，数满全放行。 机制折在下面，回查用。

⛔ 报了「三样」就三样都要有下落 ——  说「第二样跳过」等于当众漏掉一件。只念这一条就够： 「在 barrier 上等着的 warp 不腾位子 ——  所以同步密集的 kernel，占用率不会因为『大家都在等』而变好看。」

第三样值一张图。

滚到那张四行的数据通路图。只指中转那一列，别逐行念。

四行，两头一模一样 —— 左边都是全局内存，右边都是共享内存加计算单元。 在变的只有中间那一列。
最早，数据得在寄存器里过一道，地址和循环全靠线程自己算；
Ampere 起，cp.async 把寄存器那一道砍掉了，地址还是自己算；
Hopper 的 TMA，一个线程填一张描述符发出去就走 ——  地址生成整个交给硬件。

停：「走到这一步，它和 TPU 那条编译器发的 DMA 已经是同一类东西了。剩下的分歧只有一个 ——  这条 DMA 是运行时某个线程发的，还是编译期就排好的一步。」

—— 落点：记分板管「谁能发指令」，barrier 管「谁跟谁对齐」， TMA 管「数据怎么不占用线程就位」。三套加起来，就是「跑起来才决定」那笔账的明细。 TPU 一样都不需要 —— 不是它简陋，是这三个问题它在编译期就全答完了。

滚到那张「三种铺法」的大图，先只停在最上面那一栏（朴素）。 五到六分钟。这是 §3 下半段的地基，千万别赶。

两边的部件到这儿都拆完了。但有一件更基本的事，我一直没说： 一次矩阵乘，在硬件上究竟是按什么顺序算的。

⭐ 先别看图，先让台下自己说一遍。 问：「矩阵乘怎么算？」—— 十个人有十个会说： 拿左边第一行、右边第一列，对应相乘加起来，得到左上角那个数。」
让他们说出口，这一步很重要。

对。数学上完全正确。 —— 而真实硬件，无论 GPU 还是 TPU，一个都不这么算。

停。让这句话空一拍。

先数一下朴素算法这笔账。看图上第一栏。
第一步：A 的整个第 0 行乘 B 的整个第 0 列，得到 c₀₀，写出去。
第二步：还是 A 第 0 行，换 B 第 1 列 ——  注意，A 第 0 行那三个数，你又读了一遍。
四步走完：A 的六个数各读两遍，B 的六个数各读两遍，总共读 24 次， 只换来 12 次乘加。

⭐ 这里给一句好记的：「读两次，才干一次活。」 —— 而且注意：这个算法确实只要一个累加器， 台下如果有人问「那不是挺省的吗」，答案是：省错了地方。 滚到第二栏（GPU）。

硬件干的是另一件事：把 k 挪到最外层。
第一步，k 等于 0：只读两样 —— A 的第 0 列（两个数）、 B 的第 0 行（两个数），一共四个数。
然后这四个数两两配对，四次乘加， 分别加进四个格子里 —— 四个格子同时都动了。

第二步、第三步同理。三步走完，四个格子同时算完，一起倒出去。 —— 总共读 12 次，做 12 次乘加。整整好一倍。

⭐⭐ 这是整段最该被记住的一句，念慢： 「不是一个格子被反复用，是所有格子被同时慢慢填。」
顺带把那个具体问题答掉：每个格子恰好被更新 K 次， 左上和右下一模一样，没有谁多谁少。

⭐ 那为什么不能只留一行累加器？看那条算式。
累加器是 M×N 的一块，K 每走一步，读进来 M＋N 个数，能做 M×N 次乘加。 所以复用率就是 M×N 除以 M＋N。

⭐ 这条算式最狠的读法：它被短边卡死。 「你留一整行，短边就是 1 —— N 开到一百万，复用率还是 1。 长的那一边一点忙都帮不上。」
128 乘 256 的块是多少？八十五。

所以这一句是这一节真正的结论： 累加器阵列不是用来装结果的，它是用来买复用的。 你多开一个累加器，就等于让搬进来的那批操作数多配对一次。 —— 这正是第 1 节那条算术强度的主线，第一次长在了硅上。

滚到第三栏（TPU）。

TPU 换了个方向。B 先装进阵列，装完就不动。 A 一行一行从左边流进去，部分和一路往下累加 ——  从底下出来的时候，已经是算完的一行 C 了。

⭐ 这里的关键一句： 「K 方向那几项求和，是在空间上完成的 ——  三个 PE 各做一次乘加，把部分和接力往下传。这几项根本不经过外部累加器。」 滚到最下面那张对照表。这是整节的落点。

⭐ 把三栏收成一句：同一个数学，区别只在「哪个维度摊在空间上，哪个走时间」。
GPU 把 M 和 N 摊在空间上、让 K 走时间； TPU 把 K 和 N 摊在空间上、让 M 走时间。

⭐⭐ 最后这一句一定要说，它是往下三节的钥匙： 「这一个选择，决定了后面所有事 —— 累加器要多大、 tile 为什么是最核心的那个旋钮、以及 TPU 为什么非得让权重坐着不动。」

⚠️ 两条边界，讲的时候顺口带一下就行，别展开： ① 图上 GPU 那栏画的是算法级；指令级上一条 MMA 已经在内部把 K 的一小段 规约掉了 —— 形状一样，只是成批做。 ② TPU 那边外部累加器只在跨 K 分块时才用到，真机 K 常常几千， 阵列一次只吃得下 256。

这一节先别滚图。先把问题问出来，再滚 ms-g4。

上一节讲的是谁来发这条指令。这一节换一个问题 —— 一条指令发出去，一次能吃下多大一块？

听起来很内部，但它决定一件很外部的事 ——  你的 head_dim 选多少不吃亏。

现在滚 ms-g4，再切 ms-t4。两张图只看同一个位置： 求和的那一维有多深。

GPU 这边，一路五代指令，动员的线程越来越多、一次吃的块越来越大， 但求和那一维五代都没变，一直是十六。

TPU 这边，MXU 是二百五十六乘二百五十六的脉动阵列，那一维就是二百五十六。

停死，念这个数：「十六，对二百五十六。差十六倍。」 这是这一小节唯一要记住的数。

这一维有个名字，叫收缩维 ——  就是矩阵乘里被加掉的那一维，一条指令一次要沿着它累加多少项。 名字不重要，重要的是刚才那两个数。

差十六倍意味着什么？拿一个真形状套进去：注意力的 head_dim 常取 128。

GPU：一百二十八除以十六等于八，切成八条指令，一点不浪费；
TPU：一百二十八对二百五十六，只喂满一半 —— 另一半空着。

那空着的那一半，找点别的活填进去不行吗？

不行。收缩维是一条累加链 ——  塞进去的任何别的数据都会被加进同一个部分和，谁也顶不上来。

反过来问：那 GPU 为什么不吃这个亏？

不是因为它有调度器。调度器换的是 warp，填的是时间上的气泡 —— 它填不了收缩维里空着的那一半。 是因为它的 K 只有十六，几乎什么形状都是十六的整数倍。

⛔ 上面这句照着念。 早先的版本说「GPU 有调度器所以对形状不敏感」，那是错的 ——  台下懂 CUDA 的人一句话就能顶回来。

⚠️ 这不是哪边设计得差。GPU 敢把 K 做这么窄、把对齐这件事留到运行时收拾， 代价就是 3.2 里 TPU 少掉的那五样。 —— 「按谁的粒度对齐」这个问题，本来就有两个答案。

⭐ 滚到 3.4 最后那张四条带的大图（图上标题是 「四条带，从『谁不动』一路走到『要重搬几次』」） （四条带：三种驻留 ／ 切块 ／ 乒乓 ／ 重搬账）。
⛔ 别四条带平均用力，按这个优先级： ① 先讲第三条带（乒乓）—— 这是别人没讲明白的那一段。
「装一块权重要 256 个周期。但每个格子有两套权重寄存器 ——  一套在算，一套在装。所以那 256 藏在算的那 2048 里面， 只有第一块要干等。」
⭐ 然后立刻给判据：「藏得住的条件就一条：一段的行数要大于 256。」

② 再回到第一条带（三种驻留），一句话带过： 左驻、右驻、中间累加驻，区别只有一个 —— 谁不动。 TPU 走的是右驻。 而「谁不动」决定了哪两维是空间、哪一维是时间。

③ 第二、四条带留给追问： 168 块怎么来的、为什么 M 不能一口气流完（累加器要 128 MiB，比 VMEM 还大）、 段切得越碎右矩阵重搬越多次。 台下不问就不展开。

⭐ 这张图的收口只有一句，一定要说： 「段开长，右矩阵少搬；段开长，累加器变大。所以这个旋钮拧到底， 还是被片上那口灶台顶住的 —— 跟 3.6 里 GPU 那堵墙是同一堵。」 ⚠️ 这里几乎一定有人举手算账： 「一条指令才十六深，怎么也凑不出两千二百五十 TFLOPS。」
提前拦一句：指令形状是「一次交代多少活」，每周期做多少乘加是另一回事。 要对账就报课件上那条：4,096 × 2 × 148 SM × 1.83 GHz ＝ 2,218，官方 2,250，差 1.4%。 ⛔ 别顺着他那条错推法往下推。 指本小节末尾那条峰值链。这一段是给会算账的人的， 语速可以快，但结论要清楚。

顺手把 TPU 的峰值从头乘一遍，验证我们没在瞎说： 一个 MXU 一拍 262,144 次浮点运算，乘二点二 GHz，乘四个 MXU —— 等于 2,307 TFLOP/s。官方公布 2,307，一分不差。

⚠️ 这里有一条必须带走的跨代警告： TPU 那个数里有一半来自「每个 cell 每周期做两次乘加」 ——  这是二百五十六乘二百五十六这几代才有的。 一百二十八乘一百二十八的老几代，每 cell 只做一次。 跨代沿用这个常数，会直接错一倍。

这一小节的收口，一定要说： 「能融多大，由两个数一起卡死 —— 一个是工作台多大（3.2b 那个 289 倍）， 另一个就是这里的粒度多粗。3.6 切块的时候，这两个数会同时出现。」

先别滚图。这一节的问题一句话就能问完， 问完再滚 ms-p37。整节一分半，不要展开。

上半节到这儿为止，数的都是「一边有、另一边没有」。 还有两块，是各自额外长出来的。

但这一节要问的不是「它们是什么」。慢一点，问这个： 「宣传页上写着支持的东西，你到底用不用得上？」
白板上就写四个字：「用不用得上」。

这一节就两个例子，两边各一个。

现在滚 ms-p37。左边一个例子，右边一个例子， 下半张是两道门槛。

第一个例子在 GPU 这边，叫块量化。宣传口径是「B200 支持 FP4」。

这句话是真的。门槛在于 —— 带块量化的指令有两条通路， 而它们不在同一颗 die 上，B200 只有其中一条。 所以听的人默认「两条都有」，那就错了。

第二个例子在 TPU 这边，叫 SparseCore —— 一颗专门查嵌入表的核。 宣传口径是「TPU 有个专门查 embedding 表的核」。

这句话也是真的。门槛在于 —— 它只对「同一批行被反复查」的负载划算。 图上那个 duplication factor 量的就是这件事。

这里停一下，这是给这屋里的人的那一句： 「语言模型的重复度大约是一点零一 —— 一行几乎只查一次。 编译器根本不往那儿派，生产任务里它一次表都没查过。」 ⛔ 紧接着必须补一句，不能停在上面那里： 「但『用不上』不等于『没在用』—— 这颗核你天天在用。」

它在大模型训练里真正干的活是卸载集合通信 ——  把 All-Gather 和 Reduce-Scatter 从 TensorCore 手上接过去， 让 TensorCore 专心算，通信在旁边并行跑完。 这就是「算和通信重叠」在 TPU 上的落地方式之一。

⭐ 如果有人追问「不规则访存到底怎么处理的」 ——  课件里有一张 SparseCore 拆开图，台上别滚过去（会吃掉三分钟）。
口头给一句就够：「TensorCore 一次搬 (8,128) 一整块， SparseCore 一次只动 8 个数，而且地址可以是刚算出来的 ——  它分成十六个小单元各追各的地址，靠同时欠着很多次取数来扛延迟。」
要细节让他课后看那张图。 ⛔ 被追问粒度时，这两句不要说错：
「(8,128) 是向量寄存器的形状，不是 DMA 的最小粒度。」
「SparseCore 在 v7 上的 DMA 粒度是 64 字节 ——  这个数是跑 get_tpu_info 直接报出来的，不是文档里的形容词。」
还有一个陷阱问题：「它操作 VMEM 还是 HBM」——  答「两块都碰，但它碰的那块 VMEM 是它自己的，不是 TensorCore 那块」。 有人要出处就报这两个，都是公开的：Google Cloud 那篇 《Training large models on Ironwood TPUs》把它单列成一条调优手段； MaxText 公开仓库的 XLA flag 库里就有 SparseCore All Gather。 ⛔ 不要报任何内部实测数字。 ⭐ 这一节的教训到这里才完整，慢慢说： 「宣传口径不只会让你用不上它说的那个能力 ——  还会让你看不见你已经在用的那部分。你拿『查表的核』这个名字去 profile 里 找查表，找不到，就以为这颗核在闲着。」

两个例子摆完，顺带能看出一件事：这两块东西连层级都不一样 —— 块量化是一条指令的变体，SparseCore 是物理上独立的一颗核。 「两边各多出一块」这个说法本身，也是宣传口径的产物。

两句宣传语都没说错 —— 错的是把「支持」读成「我的负载用得上」。

收口给一个动作，不要只给感慨： 「看到一个『支持 X』，先去找那道门槛写在哪一行。找不到， 就当它不适用于你。」