Lecture 10: Hardware Specialization（硬件专用化）（日期：Oct 23）

目录 · ← l9 · l11 →

Lecture 10: Hardware Specialization（硬件专用化）（日期：Oct 23）

概述：本讲回答一个贯穿全课程的问题——我们前面几讲都在讲如何高效利用多核 CPU 和 GPU，那么为什么这些”通用处理器”仍然不够高效？本讲从能耗（energy）的视角重新审视计算：专用硬件（ASIC、FPGA、DSP、领域专用加速器）通过减少指令流开销、缩短数据搬运距离、用数据流（dataflow）方式组织计算，可以在每瓦性能（perf/watt）上比通用 CPU 提升 10～1000 倍。本讲以 Google TPU、NVIDIA Tensor Core、可重构数据流架构（如 Plasticine）为案例，给出”理想 AI 加速器”的特征清单，并讨论”硬件彩票”（hardware lottery）现象：硬件与算法的演化会相互塑造。

注意：本讲对应 Assignment 4（在 AWS Trainium2 加速器上编写并优化 fused Conv+MaxPool kernel），作业将让你亲身体验”软件管理片上存储 + 专用计算引擎”的编程模型，与本讲的 tiled programming、dataflow、片上数据搬运等概念直接相关。

一、核心概念与定义

1. Energy efficiency（能效）

定义：单位能量能完成的运算量。幻灯片给出关键关系： $\text{Power} = \frac{\text{Ops}}{\text{second}} \times \frac{\text{Joules}}{\text{Op}}$ 即功耗 = 吞吐率 × 每操作能耗。因此，提升能效（energy efficiency）要么提高每焦耳运算数，要么降低每运算能耗。幻灯片强调：更好的能效 ⇒ 专用化（specialization，fixed-function），并追问”专用化带来的改进幅度有多大？（What is the magnitude of improvement from specialization?）”——本讲后面的数据（10×、100–1000×）就是对这个问题的回答。
现实类比：一辆卡车（通用处理器）什么都能运但空驶率高、每吨公里耗油大；一条专用输油管道（专用硬件）只能运油，但单位运量的能耗低几个数量级。代价是管道建造成本高、只能运油。

公式/图示：

能效 (perf/watt) = Performance / Power
提高能效的两条路：
① 同样功率做更多运算（专用指令、数据流、消除指令开销）
② 同样运算花更少能量（减少数据搬运、缩短距离、低位宽数据格式）

背景动机（energy-constrained computing）：本讲开篇指出两类能耗受限场景——(1) 移动设备：电池寿命有限、无风扇散热受限（heat dissipation without fan）；(2) 超算与数据中心：机器规模巨大（数十万颗 CPU/GPU），供电与散热（cooling）都是硬约束。而 AI 的需求正以指数速度增长（AI demands are growing exponentially），数据中心严重受能耗约束——所以”能效”不是锦上添花，而是 AI 能否继续扩张的第一性约束。

2. Hardware specialization / fixed-function（硬件专用化 / 固定功能）

定义：为特定计算（如视频编码、FFT、DNN 推理）定制电路，把”程序控制”变成”固定数据通路”。相比通用处理器，专用硬件可以省掉指令取指、译码、调度等”控制开销”。
现实类比：瑞士军刀 vs 专用螺丝刀——瑞士军刀什么都能干但每样都不够趁手；专用螺丝刀拧螺丝又快又省力。
规则（来自幻灯片）：与高质量 C 代码在 CPU 上相比——
- 吞吐优化的处理器（GPU core）：约 10× perf/watt（假设代码能很好映射到宽数据并行、且为 compute-bound）；
- 固定功能 ASIC：可接近 100–1000× 甚至更高 perf/watt（假设 compute-bound 且不是浮点数学密集型）。

能效 vs 可编程性光谱（第 19 页，Pat Hanrahan 设计）：

能效 低 ←————————————————————————————→ 高
      Energy-optimized CPU (最好编程)
      可编程 DSP
      GPU (吞吐优化)        ~10×
      领域专用加速器 (TPU)   ~20×（DSL 编程，如 DNN）
      FPGA/可重构逻辑        ~50×???
      ASIC                  100–1000×（不可编程，设计费数千万美元）

注意光谱两端：DSP 最易编程；FPGA “difficult to program（making it easier is active area of research）”；ASIC “Not programmable + costs 10-100’s millions of dollars to design/verify/create”。典型固定功能场景：视频编解码、音频播放、Camera RAW 处理、神经网络（future?）。

3. Programmability overhead（可编程性开销）

定义：现代处理器执行一条指令要经过取指、译码、依赖/冒险检查、选择执行资源、读寄存器堆、数据搬移、写回、地址翻译、uop 缓存等一系列步骤。幻灯片（图源 Eric Chung）用 H.264 视频编码能耗分解说明：即使使用 SIMD，功能单元（functional units）消耗的能耗占比依然很小，大部分能耗花在寄存器读取（RF）、流水线控制（Ctrl）、流水线寄存器（Pip）、指令取指与指令缓存（IF）、数据缓存（D-$）上。
现实类比：一个 100 人的研究所，真正动手做实验的只有几个人，其余都在开会、审批、填表（控制/管理开销）。专用硬件相当于把”开会审批”全部砍掉，只留做实验的人。
量化（幻灯片，第 32 页 “Amortize overhead of instruction stream control using more complex instructions”）：
- 半精度 FMA（fused multiply-add）的”可编程性开销”约占 2000%；
- 半精度 DP4（vec4 dot product）约占 500%；
- 半精度 4×4 MMA（matrix-matrix multiply + accumulate）仅占 27%。
- 核心原则：用一条复杂指令摊销（amortize）大量运算的指令流处理成本（Key principle: amortize cost of instruction stream processing across many operations of a single complex instruction）。
复习问题（幻灯片原题）：SIMD 执行如何降低某些类型计算的开销？这些计算需要具备什么性质？——答案：SIMD 把 N 个运算合并进一条指令，指令流开销被 N 摊销；前提是运算本身可宽数据并行（相同操作、无分支）、数据布局连续。

4. DSP / VLIW（数字信号处理器 / 超长指令字）

定义：DSP 是可编程处理器，但指令流控制通路更简单；通过复杂指令（SIMD/VLIW）让每条指令完成多个运算，摊销控制成本。VLIW（very-long instruction word）：一条指令同时指定多个不同的运算（与 SIMD 的”一条指令对多份数据做同一运算”形成对比）。
现实类比：普通工人一次只能搬一件货（RISC）；VLIW 工人拿着清单一次搬五件不同的货（每个货位干什么写死在指令里），省去了逐个分配的中间管理。
实例：Qualcomm Hexagon DSP（用于 Google Pixel 手机），其 FFT 最内层循环每周期执行 29 个 RISC 级操作。

5. ASIC（Application-Specific Integrated Circuit，专用集成电路）

定义：为单一用途定制的集成电路，不可编程。代价是设计/验证/流片成本高达数千万～数亿美元。
实例 1——FFT（Chung et al. MICRO 2010）：ASIC 用约 CPU 单核 1/1000 的芯片面积、约 1/100 的功耗 达到同样性能；GPU core 比 CPU core 面积效率高约 5–7×。
实例 2——Anton 超算（DE Shaw Research）：Anton 1（2008）专为分子动力学蛋白质模拟设计，512 个 ASIC 计算粒子-粒子相互作用，配吞吐优化的 FFT 子系统与为 N-body 通信模式定制的低延迟网络；Anton 3（2025）比同时代 GPU 快约 20 倍。
现实类比：专线大巴（ASIC）与出租调度中心（通用 CPU）——大巴线路固定但每公里成本低得多，前提是客流（工作负载）足够稳定。

6. FPGA（Field Programmable Gate Array，现场可编程门阵列）

定义：ASIC 与处理器之间的”中间地带”：芯片上排列逻辑块（logic blocks），通过可编程互连连接。程序员定义的逻辑直接由 FPGA 实现，无需流片。
LUT 机制：可编程查找表。Xilinx Virtex-7 的 LUT6 是 6 输入 1 输出，可视为 64 元素真值表；例如 6 输入 AND 用 1 个 LUT6；40 输入 AND 通过级联 8 个 LUT6 实现（延迟 = 3）。
现代 FPGA：大量面积用于”硬门”（hard gates）——SRAM 存储块、DSP 块（乘法器）、CPU（ARM、RISC-V）；用硬件描述语言（Verilog 等）编程。AWS EC2 F1/F2 提供云端 FPGA。
能效定位：约 50×？（幻灯片标注”jury still out”，尚无定论），编程困难（如何让它更好编程是活跃研究领域）。
现实类比：乐高积木（FPGA）——你可以随时拆了重搭成任何结构（ASIC 的能力），但每个积木块本身比定制浇筑的钢筋混凝土（ASIC）低效。

7. Domain-specific accelerator（领域专用加速器）

定义：介于”完全可编程”与”完全固定”之间：在有限领域内可编程，通常通过领域专用语言（DSL，如 DNN 场景）编程。约 20× 能效（幻灯片以 Google TPU 为例）。
现实类比：食堂的”套餐窗口”——菜品组合固定几种（领域受限），但比”点菜窗口”（通用）出餐快得多，又比”中央厨房专线”（ASIC）灵活。
关键词：TPU 的关键指令（v1 幻灯片）：read host memory、write host memory、read weights、matrix_multiply / convolve、activate——指令集小到只有几个操作。
TPU v1 的芯片面积分配（第 51 页，Jouppi et al. 2017）：算术单元约占芯片的 30%，控制逻辑占面积很小（”Note low area footprint of control”）——对比通用 CPU 中大量面积被控制/缓存占据，这正是”专用化省面积”的直接证据。
GPU 为什么对 AI 不是最优平台？（第 22 页的提示问题）——AI 模型计算特征接近稠密矩阵乘（high arithmetic intensity），GPU 有丰富 FLOPs 与 cuDNN 等优化库，但 GPU 本质仍是通用可编程处理器：它为任意数据并行程序保留了指令流、调度、缓存等机制，而这些对 AI 工作负载是”多余的”（hint: is a general purpose processor needed?）。专用硬件可以把这些开销全部去掉，把面积和功耗让给算术单元。

8. Systolic array（脉动阵列）

定义：由大量同构处理单元（PE）排成阵列，数据像”脉搏”一样在相邻 PE 之间逐周期流动；权重驻留在各 PE（或从 FIFO 流入），输入/部分和沿阵列传播。这是数据驱动（data-driven，wavefront）执行，而非指令驱动。
现实类比：工厂流水线上的工人传递半成品——每个工人（PE）只做自己那一小步（乘加），把结果传给下一个工位，原料从一端流入，成品从另一端流出；没有人需要”下令”每个工位做什么（固定功能），只需让物料流动。

SIMD vs Systolic array 对比（幻灯片表格）：

特征	SIMD	Systolic Array
Dataflow	Control-driven（指令驱动）	Data-driven（wavefront 数据驱动）
Locality（数据重用）	Limited	Temporal + Spatial（时间与空间）
Communication	Global（寄存器/内存）	Local（相邻 PE）
Control	Centralized（集中控制）	Distributed（分布式）
Efficiency（perf/mm², perf/Watt）	Medium	Very high

如何构建更大的矩阵乘（第 60–63 页）：实际 GEMM 远大于 4×4（例：A=8×8，B=8×4096，C=8×4096）。做法：让脉动阵列的累加器（assume 4096 accumulators）分时处理不同的输出列块——把 B 的 4096 列切分成列块，权重（A 的行/列）按块复用，阵列反复执行小规模 MMA，把结果累加到 4096 个累加器中。即：片上阵列做”原子 MMA”，大 GEMM 由阵列+累加器+分块调度组合完成——这也是后续 tensor core（H100 989 TFLOPS）与 tiled programming（16×16/32×32 tile）的雏形。

9. Dataflow / Streaming dataflow（数据流 / 流式数据流）

定义：把计算组织成数据流图：AI 模型本身就是一个 dataflow graph（Weights → GEMM1 → Pool → GEMM2 → SoftMax → Sum …）。专用硬件直接把中间数据在计算单元之间流动，用”在芯片上按空间布局调度计算”代替逐条取指执行。理想情况下没有指令 ⇒ 没有取指/译码开销；极端异步：没有顺序指令执行。
现实类比：水流过一连串水轮（每个水轮做固定工作）发电，而不是”一个机器人反复读说明书操作同一个水轮”。水轮之间用管道（数据通路）直连，中间产物不落地。
关键特性（理想加速器清单）：tiled tensors（16×16、32×32）实现 GEMM 最大 TFLOPS 与低指令开销；异步计算、异步访存、异步 chip-to-chip 通信（重叠 compute/memory/communication）；compute-unit 到 compute-unit 的通信；fusion 与 pipelining；streaming dataflow；避免 off-chip 数据访问。

10. Tensor core + TMA + Tiled programming（张量核 / 张量内存加速器 / 分块编程）

定义：GPU 上的专用矩阵乘单元（tensor core）执行一条 MMA 指令完成整个矩阵分块乘加；TMA（Tensor Memory Accelerator，H100）用专用指令异步搬运张量分块；tiled programming（CUTLASS、Triton、Thunderkittens）以 16×16、32×32 的 tile 为编程原语，目标是 GEMM 最大 TFLOPS、低指令开销。
关键事实（幻灯片）：
- A100 的每个 SM：64 个 fp32 mul-add ALU、32 个 int32 ALU、4 个 tensor core（执行 8×4 × 4×8 矩阵乘加，A、B 为 fp16，D 为 fp32 累加）；GA100 共 108 个 SM → 19.5 TFLOPs fp32 + 312 TFLOPs fp16/32 mixed（tensor cores）。
- H100（2022，TSMC 4nm，800 亿晶体管）：第四代 tensor core、TMA、CUDA cluster、HBM3 最高 80GB；144 个 SM；tensor core（systolic array MMA）989 TFLOPS fp16，SIMD 134 TFLOPS fp16 / 67 TFLOPS fp32。
- “All the TFLOPS are in the Tensor Cores”：幻灯片显示各代 GPU 的浮点算力中 tensor core 占比 89%、50%（或 0%？）、94%、96%、98%——绝大多数算力来自张量核。
H100 的分层结构（第 38–39 页）：CUDA 层级 Grid → Cluster → Thread Block → Threads/SIMD Lanes，对应计算层级 GPU → CPC → SM → SIMD Lanes，对应存储层级 80 GB HBM/50 MB L2 → 256 KB shared memory/SM → 256 KB shared memory → 1 KB RF/thread、64 KB/SM partition。Thread block cluster 是最多 16 个 thread block 的集合，保证每个 thread block 在不同 SM 上同时执行（支持跨 SM 协作）。每个 SM 内：4 个 warp scheduler（每周期各取指/译码 1 个 warp），64 KB 寄存器/sub-core、256 KB 寄存器/SM，按最多 64 个 warp 划分；含 SIMD fp32/int/fp64 单元、LSU、4 个 16×16×16 [fp16 fp16 fp32] tensor core 单元。
TMA 细节（第 40 页）：专用指令高效搬运数据；异步地把张量的一个区域从 global 加载/存储到 shared memory；拷贝描述符（copy descriptor）描述区域；单个线程发起 TMA 操作（cuda::memcpy_async）；拷贝完成时信号屏障；硬件完成地址生成与数据搬运——程序员不需要逐元素寻址。
B100 的变化（”Not your father’s CUDA”）：寄存器带宽限制 tensor core；张量数据放 SMEM 与 TMEM；单个线程执行 MMA ⇒ 不再有 warp 的概念；编程步骤：tcgen05.alloc 分配 TMEM 与描述符 → cp.async.bulk.tensor 用 TMA 预取/流式搬运 tile（配合 mbarrier）→ tcgen05.mma batch 启动异步 MMA（配合 tcgen05.commit）→ tcgen05.fence 排序与回收。
现实类比：tile 编程就像”集装箱运输”——不再逐件搬运（标量）或逐托盘搬运（SIMD），而是整个集装箱（16×16 tile）用专用吊机（tensor core + TMA）一次吊装。

11. Numerical formats（数值格式：BF16 / BF8）

定义：为 AI 计算设计的低精度浮点格式，降低每运算能耗与面积。
- BF16：1 符号位 + 8 指数位 + 7 尾数位。与 FP32 相同范围（range），但精度更低。
- BF8 E4M3：1+4+3，最大值 448。
- BF8 E5M2：1+5+2，最大值 57344。
现实类比：记账用”元”（FP32）需要更多位数；估预算用”万元”（BF16）范围一样大、但精度粗——对神经网络这种”对噪声不敏感”的计算足够，且省电省面积。
公式：FP32 表示 ≈ $-1^S \times (1 + M \times 2^{-23}) \times 2^{E-127}$；BF16 把尾数位砍到 7 位，保留 8 位指数（范围同 FP32）。

12. Hardware lottery（硬件彩票）

定义（Sara Hooker 定义）：当一个研究想法获胜，是因为它恰好适合当时可用的软件与硬件，而不是因为它普遍优于其他研究方向时，就发生了”硬件彩票”现象。
幻灯片图示：TPU 擅长 dense matrix multiply（运算强度 OI ∝ n）→ 人们设计 Transformer 模型 → Transformer 模型主导 → 硬件进一步为 MM 专用化。硬件与算法形成正反馈循环。
现实类比：中彩票者把成功归因于自己的”策略”，但其实只是抽到了适合当时机器的号码；当机器变了，同一策略可能一文不值。
延伸解读：这个循环也解释了第 48 页的观察——”AI is the driving force behind new architectures, compilers, and system design”（AI 正在重新定义计算）：TPU3、Apple Neural Engine、AWS Trainium 2、Cerebras Wafer Scale Engine、SambaNova 等纷纷为 AI 造芯。算法研究者与硬件设计师需要意识到：当前架构偏好可能只是历史偶然，而不是最优解——这也正是课程鼓励学生思考”通用 vs 专用”长期权衡的原因。

13. 补充：理想 AI 模型加速器的特征清单（贯穿全讲的 checklist）

幻灯片第 23–29 页反复出现一张”特征-原因”表，作为评估任何加速器（GPU、TPU、数据流芯片）的标尺，在此汇总：

特征（Feature）	原因（Why?）
Tiled tensors（16×16、32×32）	GEMM 最大 TFLOPS、低指令开销
异步计算（asynchronous compute）	重叠计算与访存
异步内存访问（asynchronous memory access）	重叠计算与访存
异步 chip-to-chip 通信	重叠计算、访存与通信
计算单元间通信（compute-unit-to-compute-unit comm.）	数据在片上直接流动
Fusion 与 Pipelining	减少中间结果落盘
Streaming dataflow	数据流式执行、无指令开销

第 47 页用这张表”打分”了 NVIDIA GPU：Tiled tensors ✅（CUTLASS/Triton/TK 生态）、异步计算 ✅（mma_async）、异步访存 ✅（TMA+TMEM）、异步芯片间通信 ❓（TB Cluster，仍在探索）——说明即使是最先进的 GPU，离”理想加速器”也还有一步之遥，而这正是数据流架构（第 68–70 页，Plasticine/SambaNova）试图补齐的部分。

二、代码示例与详细解说（本讲重点）

示例 1：脉动阵列的 PE 级伪代码（矩阵-向量乘 y = W·x）

// 伪代码：4x4 脉动阵列执行 y = W·x（W 为 4x4 权重矩阵）
// 硬件结构：
//   - 16 个 PE 排成 4x4 网格
//   - 权重 w_ij 预先装入 PE(i,j)（或由 Weights FIFO 逐周期供给）
//   - 输入向量 x 从左侧逐元素流入，向右传播（每个周期前移一格）
//   - 部分和从上往下传播，最后一行 PE 输出到 32-bit accumulator
//   - 每个 PE 内部：1 个乘法器 + 1 个加法器 + 1 个局部寄存器

// 每个 PE(i, j) 每周期执行的逻辑：
loop forever:
    x_in   = receive_from_left(PE(i, j-1))   // 从左边邻居拿输入
    acc_in = receive_from_top(PE(i-1, j))    // 从上方邻居拿部分和
    acc_out = acc_in + w[i][j] * x_in        // 乘加：本 PE 的唯一运算
    send_to_right(PE(i, j+1), x_in)          // 输入继续向右流动
    send_to_bottom(PE(i+1, j), acc_out)      // 部分和向下流动

// 时序（wavefront）：第 t 周期，x0 到达 PE(0,0)；
// t=0: x0*w00 在 PE(0,0) 累加
// t=1: x0 传到 PE(0,1)，x1 进入 PE(0,0)；PE(0,0) 把部分和 x0*w00 传给 PE(1,0)
// t=2: PE(0,2) 得 x0*w02，PE(1,1) 得 x0*w10+x1*w11，PE(2,0) 得 x0*w20 …
// 最终：PE(3,j) 依次输出 y_j = Σ_i w_ij * x_i（对应列累加）

数据流动示意（wavefront 的三次快照）：

t=1:                               t=2:
┌────┬────┬────┬────┐              ┌────┬────┬────┬────┐
│x0*w00→ x0 → x0 → x0│             │x1*w00→x0*w01→ x0 → x0│
│  ↓                          │             │  ↓                            │
│x0*w00→  0 →  0 →  0 │             │x1*w10→x0*w11→  0 →  0 │
│  ↓                          │             │  ↓                            │
│  0    0    0    0  │             │x1*w20→  0    0    0  │
│  0    0    0    0  │             │  0    0    0    0  │
└────┴────┴────┴────┘              └────┴────┴────┴────┘
（x0 向右传播，部分和向下传播）        （x1 进入，部分和继续下移）

【代码做了什么？】 这段伪代码描述了一个 4×4 脉动阵列如何计算矩阵-向量乘 $y = Wx$。所有 16 个 PE 每一拍（cycle）同时工作：每个 PE 把从左邻收到的输入 $x_i$ 与本地权重 $w_{ij}$ 相乘，加到从上邻收到的部分和上，然后把输入继续向右传、把部分和继续向下传。权重事先固定（或者通过 Weights FIFO 逐拍注入），整个计算没有”取指令、译码、判断分支”等任何控制流操作——PE 内部只有一条固定的乘加数据通路。幻灯片用多张图逐拍展示了 $x_0 w_{00}$、$x_0 w_{00}+x_1 w_{01}$、$x_0 w_{00}+x_1 w_{01}+x_2 w_{02}$ 等部分和在阵列中”涌动”的过程，这就是 wavefront（波前）的含义。

【并行机制解说】 这里的并行性是空间并行（spatial parallelism）+ 时间流水（temporal pipelining）的结合：16 个 PE 同时做不同的乘加（空间），同时数据像流水线一样逐拍推进（时间）。与 SIMD 的根本区别在于通信模式：SIMD 的每条指令都要访问全局寄存器/内存（global communication，受寄存器带宽限制），而脉动阵列只做相邻 PE 之间的本地通信（local，neighbor PEs）——数据在芯片上”走最短的路”，不需要经过集中式寄存器堆。这对应本讲概念 8（systolic array）与概念 9（dataflow）：控制是分布式的（每个 PE 没有自己的指令流），能效极高（perf/mm²、perf/Watt 都是 “very high”）。矩阵-矩阵乘（Y=WX）时，需要多组 4×32-bit accumulator 来容纳输出列（幻灯片第 58 页的”Notice: need multiple 4x32bit accumulators to hold output columns”）。

示例 2：能耗对比——为什么专用硬件省电（减少指令取指开销）

// 能耗数量级对比（幻灯片 [Dally / Olson] 的 ballpark 数据）
// 单位：皮焦耳 pJ（10^-12 焦耳）
//
//  整数运算 (integer op)                    ≈ 1 pJ
//  浮点运算 (floating point op)             ≈ 20 pJ
//  读 64 bit 本地小 SRAM（片内 1mm 远）      ≈ 26 pJ
//  读 64 bit 低功耗移动 DRAM (LPDDR)        ≈ 1200 pJ
//
// 注：以上仅为"做逻辑运算本身"的能耗，不含指令译码、寄存器装载等开销

能耗数量级速查表（设计决策时随身携带）：

操作	能耗	相对整数运算
整数运算	~1 pJ	1×
浮点运算	~20 pJ	20×
读 64 bit 片内 SRAM（1mm）	~26 pJ	26×
读 64 bit LPDDR	~1200 pJ	1200×

// 反例推演：在通用 CPU 上算 y = W·x（4x4）
// 每条标量指令都要：取指(IF) + 译码 + 冒险检查 + 读寄存器堆(RF) + 搬数据 + 运算 + 写回
// H.264 编码实测能耗分解 [Hameed ISCA 2010]：功能单元(FU)占比很小，
// 大头是 RF、Ctrl、Pip、IF、D-$
// 专用加速器做法：把"指令开销"全部砍掉
//   FMA (半精度标量乘加)   → 可编程性开销 ≈ 2000%   （指令开销是运算的 20 倍）
//   DP4 (半精度 vec4 点积) → 可编程性开销 ≈ 500%
//   4x4 MMA (矩阵乘加)     → 可编程性开销 ≈ 27%

【代码做了什么？】 这个”示例”其实是一组设计准则与数量级数据，用来回答”为什么通用处理器效率低、专用硬件为什么省电”。幻灯片给出两条主线：(1) 执行一条指令本身很贵——现代处理器执行一条指令要读指令、译码、检查依赖/冒险、找执行资源、控制寄存器堆 SRAM、搬数据、运算、写回、转 uop、访问 uop 缓存、地址翻译、访问 icache 等（第 8 页）；H.264 编码的能耗分解显示功能单元消耗的占比很小，绝大部分能耗在”伺候指令”上。(2) 数据搬运比运算贵得多——读一次 LPDDR（1200 pJ）相当于做 1200 次整数运算（1 pJ）或 60 次浮点运算（20 pJ）。所以设计系统时第一原则永远是”减少数据搬运”。

【并行机制解说】 这个示例解释了两个能效来源：(a) 用复杂指令摊销控制成本——一条 MMA 指令完成 64 次乘加，指令流处理成本被摊销到 64 个运算上，可编程性开销从 2000%（标量 FMA）降到 27%；这正是 SIMD 的延伸思路（回顾课上问题：SIMD 如何降低开销？要求运算可宽数据并行、无分支）。(b) 固定功能消除控制——ASIC/脉动阵列根本没有指令流，能效可达 CPU 的 100–1000×（对应概念 2、3、8）。而”减少数据搬运”则引出 tiled programming 与数据流架构（概念 9、10）：数据尽量留在片上 SRAM、在计算单元间直接流动，而不是反复进出 DRAM。这也是本讲后面”理想 AI 加速器”特征清单的由来。

示例 3：Tiled GEMM 的调度伪代码（数据流加速器 / GPU tensor core）

// 伪代码：以 16x16 / 32x32 tile 为单位调度 GEMM: C[M,N] = A[M,K] * B[K,N]
// 编程模型对应：CUTLASS / Triton / Thunderkittens 的 tiled 抽象
// 硬件目标：最大化 tensor core 利用率，最小化数据搬运

GEMM_TILED(A, B, C, M, N, K):
    # 外层循环：遍历输出 tile
    for m_tile in range(0, M, 16):          # 输出行方向分块
        for n_tile in range(0, N, 16):      # 输出列方向分块
            acc[16][16] = 0                 # 片上累加器（寄存器/SMEM）
            # 内层循环：沿 K 方向累加
            for k_tile in range(0, K, 32):
                # 异步搬运：把 A、B 的 tile 从 global memory 搬到 SMEM
                a_tile = async_load(A[m_tile:m_tile+16, k_tile:k_tile+32])   # TMA
                b_tile = async_load(B[k_tile:k_tile+32, n_tile:n_tile+16])   # TMA
                wait_tiles_ready(a_tile, b_tile)            # mbarrier 同步
                # 异步计算：tensor core 执行 MMA，与下一次搬运重叠
                mma_async(acc, a_tile, b_tile)              # 16x16x32 MMA
            store_tile(C[m_tile:m_tile+16, n_tile:n_tile+16], acc)   # 异步写回

【代码做了什么？】 这是”理想 AI 加速器”编程模型的骨架：把大 GEMM 切成 16×16（或 32×32）的 tile，两层循环遍历所有输出 tile，内层沿 K 方向累加。关键动作有三个：异步加载（把 A、B 的 tile 用 TMA 之类的专用单元搬进片上 SRAM，同时干别的）、屏障同步（mbarrier 等 tile 就绪）、异步 MMA（tensor core 执行矩阵乘加，同时下一批 tile 已经在搬运）。最后把累加结果写回 HBM。幻灯片用三张”理想加速器”表格强调：tiled tensors 的目的就是 GEMM 最大 TFLOPS + 低指令开销；异步 compute / 异步 memory / 异步 chip-to-chip communication 的目的都是重叠计算与访存。

【并行机制解说】 这里的并行来自三个层面：(1) 数据并行：不同输出 tile 之间完全独立，可由不同 SM/线程块并行计算；(2) 流水线并行：同一 tile 的”搬运下一块、计算当前块、写回上一块”三个阶段通过 asynchrony 重叠——这正是”异步（非阻塞）执行”（第 24 页：Start later operations before earlier operations are complete），避免访存时计算单元空转；(3) SIMD 级并行：tensor core 本身是 systolic-array 形式的 MMA 单元（H100 的 989 TFLOPS fp16 全部来自这些单元）。软件侧（CUTLASS/Triton/Thunderkittens）负责把 tile 布局、流水阶段、屏障精确映射到硬件；硬件侧（TMA + mbarrier + async MMA）提供低开销的异步原语。这对应本讲概念 9、10——”GEMM 计算便宜，但数据搬运贵（面积、瓦、纳秒）”，所以整个调度的核心目标是让数据尽量在片上流动、让 tensor core 永不空闲。

示例 4：AI 模型 = 数据流图 → 可重构数据流处理器映射（Plasticine）

// 伪代码：把 AI 推理模型（数据流图）映射到可重构数据流架构
// 硬件：Plasticine [Prabhakar, Zhang et al. ISCA 2017]
//   PCU (Pattern Compute Unit)：可配置的计算单元
//   PMU (Pattern Memory Unit)：可配置的存储单元
//   S  (Switch)：片上互连开关

// 模型（以 CNN 为例）：
//   Sample → GEMM1(卷积/全连接) → Pool → GEMM2 → SoftMax → Sum → 输出

// 映射方式：把图中每个算子布局到芯片上的不同 PCU/PMU，
// 数据沿物理连线“流”过整条链，而不是每步取指令执行
place(GEMM1,  pcu[0..3]);    # 4 个 PCU 做 GEMM1 的分块计算
place(Pool,   pcu[4]);       # 1 个 PCU 做池化
place(GEMM2,  pcu[5..8]);    # 4 个 PCU 做 GEMM2
place(SoftMax,pcu[9]);       # 1 个 PCU 做 softmax
route(S, GEMM1_out -> Pool_in);    # 开关把中间结果直接送进下一级
route(S, Pool_out  -> GEMM2_in);
route(S, GEMM2_out -> SoftMax_in);
# 运行时：无指令流、无全局时钟同步，
# 数据到达即计算（token 控制，无需 lock 同步）

# 更细粒度：FlashAttention 类算子可按 tile 流水化
# Tile0..Tile15：QK^T → Mask → Softmax → Dropout → ×V，
# 每个 tile 在独立的 PCU/PMU 链上流式执行（kernel fusion 的效果）

【代码做了什么？】 这段伪代码描述”AI 模型 ⇒ 数据流处理器”的映射思想（第 68 页）：因为 AI 模型本身是数据流图（GEMM + 并行模式 map/filter/reduce 等），所以可以用可重构数据流架构（Plasticine 是研究原型，SambaNova 是商用版）把图的每个节点”铺”在芯片的不同计算/存储单元上，用开关网络把中间数据在单元间直接流动。运行时没有指令取指/译码，没有顺序指令执行，数据到达即计算（第 69 页：”No instructions ⇒ No instruction fetch/decode overhead；Extreme asynchrony: no sequential instruction execution”）。下半部分展示 kernel fusion 的 tile 级流水：FlashAttention 的 QK^T、Mask、Softmax、Dropout、×V 按 tile 流式执行，中间结果永远不落回 off-chip 内存。

【并行机制解说】 这是”空间调度计算（schedule computation by laying it out spatially on the chip）”的极致体现（第 71 页总结）：(1) 不同算子在不同 PCU 上空间并行；(2) 同一算子的不同 tile 在不同 PCU/PCU 链上流水并行（metapipeline/粗粒度流水，第 70 页）；(3) 数据在单元间直接流动，避免经过 DRAM，从而规避”读 LPDDR 1200 pJ vs 片内 SRAM 26 pJ”的巨大能耗差（概念 9、11 结合）。与 GPU 相比，数据流架构没有 kernel launch 开销、没有 lock 同步，只有 token 级数据流控制。这对应概念 9（dataflow）、概念 2（specialization）、概念 3（无指令开销）——也是”理想加速器”清单最后两项（compute-unit-to-compute-unit comm.、fusion & pipelining、streaming dataflow）的实现方式。

FlashAttention 的 tile 级数据流（第 70 页）：把 attention 的计算（QK^T → Mask → Softmax → Dropout → ×V）按 tile 切分（Tile 0..15），每个 tile 在独立的 PMU/PCU 链上流式执行：Q 与 K^T 的 GEMM 在 PCU 上算，Mask/Softmax/Dropout 各占一个 PMU/PCU 级，结果直接流入 ×V 的下一级。16 个 tile 的处理形成一条 MetaPipeline（元流水线）——这是本讲末尾埋下的伏笔，第 11 讲将详细展开 metapipelining 的编程模型。

三、关键要点

能效是新的性能指标：Power = Ops/s × Joules/Op。移动设备、超算、数据中心、AI 全部受能耗约束；AI 需求指数增长而数据中心能耗受限，”每焦耳运算数”比”峰值 FLOPS”更能决定系统价值。
通用处理器的低效主要来自控制与数据搬运，而不是运算本身：H.264 能耗分解中功能单元占比很小；读一次 LPDDR（1200 pJ）≈ 1200 次整数运算（1 pJ）。结论：减少指令流开销 + 减少数据搬运 = 能效提升的两大杠杆。
能效梯度（相对高质量 C 代码）：GPU ~10×；领域专用加速器（TPU）~20×；FPGA ~50×（未定论）；ASIC 100–1000×。可编程性与能效此消彼长：programmability adds overhead ⇒ reduces efficiency。
脉动阵列 / 数据流 = 专用硬件的核心组织方式：数据驱动（wavefront）、相邻 PE 本地通信、分布式控制、时间+空间双重数据重用，perf/mm² 与 perf/Watt 都远高于 SIMD。
GPU 正在走向专用化：tensor core 占了绝大多数 TFLOPS（H100 989/1123 ≈ 88%+），并引入 TMA、TMEM、tcgen05、FP8/FP4、Transformer Engine 等专用机制；”All the TFLOPS are in the Tensor Cores”，程序员必须学会用 tile 级、异步的编程模型（CUTLASS/Triton/Thunderkittens）才能榨出性能。
硬件彩票提醒我们硬件与算法互相塑造：TPU 擅长 dense MM（OI ∝ n）塑造了 Transformer 的统治地位，Transformer 又反过来推动硬件进一步为 MM 专用化。
从”为什么专用”到”怎么编程专用”：本讲结尾把问题抛给下一讲——H100/B100 的异步机制（TMA、TMEM、tcgen05）让编程变得极其复杂（”Not your father’s CUDA”），需要 ThunderKittens 等 DSL；数据流架构（Plasticine、SambaNova）则用”无指令 + metapipelining”提供更简单的编程模型。

四、常见陷阱与注意事项

只算峰值性能、忽略能耗：幻灯片反复强调 FFT 案例——ASIC 用 1/1000 面积、1/100 功率达到同等性能。真实系统中（数据中心、手机）能耗预算才是硬约束；”更快但更耗电”的方案往往不可行。注意 ASIC 100–1000× 的前提是 compute-bound 且非浮点数学密集型，不要套用到所有负载。
忽略指令流开销：在通用处理器上写标量循环时，90% 以上的能耗可能花在取指/译码/寄存器堆上而不是运算上。优化时不要只盯着 ALU 利用率，还要看指令吞吐与开销摊销（用 FMA/DP4/MMA 这类复合指令摊销控制成本）。
忽视数据搬运（memory wall）：读 LPDDR 比做整数运算贵三个数量级。即使算法运算量最优，如果数据反复进出 DRAM，能效和性能都会崩塌。设计加速器/ kernel 的第一原则是”减少 off-chip 数据访问”（tiling、fusion、流式数据流）。
加速器利用率不足：专用硬件只对匹配的负载高效。把不规则、低运算强度（arithmetic intensity）的代码硬塞给 tensor core / 脉动阵列，利用率可能远低于峰值——”assuming code maps well to wide data-parallel execution and is compute bound”这个前提不能丢。
低估专用化的成本与风险（hardware lottery）：ASIC 设计/验证/流片要数千万～数亿美元，且押注的算法可能被新研究取代（Transformer 时代硬件跟着 MM 走）。FPGA/可重构/DSA 是”可编程性与能效的折中”，但编程难度是活跃研究问题，不要以为”专用 = 免费”。
把数值格式当 FP32 用：BF16 与 FP32 范围相同但精度低，BF8 精度更低；低精度格式省能耗与面积，但需要理解精度-能耗权衡，训练/推理中要做精度管理（如混合精度累加 fp32）。
忽视”面积效率”这个维度：FFT 案例中 ASIC 用 1/1000 芯片面积达到同等性能——面积小意味着同样大小的晶圆能产出更多芯片、成本更低。评估专用硬件时 perf/mm² 与 perf/watt 同样重要。

五、思考题（带答案）

Q1：幻灯片给出 H.264 编码中功能单元能耗占比很小的数据。请用 Power = Ops/s × Joules/Op 解释：为什么”用 SIMD 提升每指令运算数”能提升能效？这种提升有什么前提条件？

答案：SIMD 把 N 个标量运算合并到一条指令里，指令流（取指、译码、控制）的能耗只付一次，摊到 N 个运算上，即”每操作能耗 Joules/Op”下降，而 Ops/s 不变（或上升），所以 Power 上升但每运算能耗下降，perf/watt 提升。前提是：运算本身可宽数据并行（相同操作、无分支）、且是 compute-bound（运算时间占主导，指令开销占比大才有得省）。这正是课堂上”SIMD 如何降低开销”的复习问题。

Q2：为什么脉动阵列（systolic array）比 SIMD 能效更高？请从数据流、通信、控制三个维度对比，并说明这对”数据搬运能耗”（26 pJ SRAM vs 1200 pJ LPDDR）意味着什么。

答案：三个维度：(1) 数据流——SIMD 控制驱动（每条指令从寄存器堆取数据），脉动阵列数据驱动（wavefront，数据自己流动）；(2) 通信——SIMD 走全局寄存器/内存（集中式，受寄存器带宽限制），脉动阵列只走相邻 PE（本地、分布式）；(3) 控制——SIMD 集中控制（指令流开销大），脉动阵列分布式/无控制（固定功能）。因为数据只在相邻 PE 间移动、且可时间+空间重用，数据搬运距离短、次数少，绕开了”读 DRAM 1200 pJ”的天价能耗，主要花在便宜的片上 SRAM（26 pJ）甚至寄存器级流动上。

Q3：假设你要为一家公司设计一个”理想 AI 模型加速器”。请根据幻灯片给出的特征清单（tiled tensors、异步 compute/memory/communication、fusion & pipelining、streaming dataflow），说明为什么这五条能同时提升”compute-bound 模型的性能上界”与”BW-bound 模型的性能上界”。

答案：tiled tensors（16×16/32×32）保证 GEMM 达到最大 TFLOPS 且指令开销低——直接决定 compute-bound 模型（如大矩阵乘）能贴住算力 roofline；高内存带宽 + 异步 memory access 让 BW-bound 模型（如逐元素/attention 类）能贴住带宽 roofline；异步 compute 与异步 chip-to-chip communication 把”访存、通信、计算”三段重叠，使任何一段都不成为短板；fusion & pipelining + streaming dataflow 把中间结果留在片上流动，既降低 off-chip 数据访问（省能耗、提带宽利用率），又让多个算子的执行流水重叠。五条合起来就是：让计算永远不空闲、让数据尽量不离开芯片、让搬运与计算重叠——这正是一个加速器同时逼近计算与带宽两个性能上界的充分条件。

Share on

Bluesky Facebook LinkedIn Mastodon X (formerly Twitter)

Runguo Li