Lecture 18: Heterogeneous Parallelism, Hardware Specialization

目录 · ← l17 · l19 →

Lecture 18: Heterogeneous Parallelism, Hardware Specialization

1. 章节标题与概述

Lecture 18: Heterogeneous Parallelism, Hardware Specialization（异构并行与硬件专用化）

• 本讲核心问题：当”把芯片做快”这件事在功耗墙上撞死之后，多出来的晶体管该拿去做什么？ 前半讲从一条最小的观察出发——一个优化良好的并行实现能比单线程 gcc -O3 的 C 代码快约 44 倍——追问”如果让你买一台新机器，你会选 4 个各性能为 P 的核，还是 16 个各性能为 P/2 的核”。答案取决于程序里可并行部分的比例，而这正是 Hill & Marty 的 “Amdahl’s Law in the Multicore Era” 要形式化的问题。后半讲把问题推到极致：Dennard 缩放（等比例功耗缩放）在 ~2003 年终结，芯片功耗必须恒定，于是每一代新晶体管里有约 30% 根本”点不亮”（暗硅 dark silicon）——结论是必须用专用化来换能效：从延迟优化的通用核，到吞吐优化的 GPU 核（~10× perf/W），到固定功能 ASIC（~100–1000× perf/W），再到 FPGA 与粗粒度可重构阵列（CGRA）。

• 涉及的主要硬件/软件机制：硬件侧包括同构多核 / 非对称（asymmetric）异构多核 / 大小核（big.LITTLE、Alder Lake）/ SoC 上的固定功能单元（视频编解码、Neural Engine、ISP、加密）/ 可编程 DSP（Qualcomm Hexagon，VLIW）/ FPGA（LUT6、硬核 DSP 与 BRAM）/ ASIC（Google TPU、Anton 的粒子相互作用单元）/ 脉动阵列（systolic array）与张量核（tensor core）；软件侧包括主机 + 加速器的异步任务模型（LD/ST/AO 重叠、TMA 张量搬运、TMEM、mbarrier）、CUDA 的 tiled-tensor 编程模型（CUTLASS / Triton / Thunderkittens）、数据流图与算子融合（FlashAttention、MetaPipeline）以及领域专用语言（DSL）。

• 在并行计算知识体系中的角色：本讲是整门课的”收束与转向“。前面十几讲把同构多核/GPU 上的正确性（一致性、同步、无锁）与性能（work-span、带宽、局部性、Roofline）讲透，其隐含前提始终是”所有处理器都能跑所有任务，只要让所有处理器一直忙”。本讲打破这个前提：目标是”把每个任务放到最合适的那类单元上”，而不是”让所有单元都忙”。它把能效（perf/Watt）提升为与性能并列的第一等设计目标，并把调度与算法分解的难题推给程序员——这与课程后续关于领域专用加速器（DNN 训练/推理、图计算）以及系统设计的讨论直接衔接。

• 配套材料：

  • CMU 15-418/618 Fall 2026 本讲讲义：已公开（可公开下载）。 Fall 2026 日程表 https://www.cs.cmu.edu/~418/schedule.html 中 Lecture 18 一行的标题为 “Heterogeneous Parallelism, Hardware Specialization”，日期为 Oct 7，其 slides 链接指向两份 PDF：lectures/20_heterogeneity.pdf（44 页，”Heterogeneous Parallelism and Hardware Specialization”）与 lectures/20_specialization_csd.pdf（53 页，”The inexorable rise of hardware specialization”），均位于公开目录 https://www.cs.cmu.edu/~418/lectures/ 之下。两点说明：①这两份 PDF 的文件名编号是 20（沿用历史学期的讲次编号），首页写有 “CMU 15-418/15-618, Spring 2025” 或页脚标注 “15-418/618 Fall ‘25”，这是讲义沿用的正常现象，不视为错误；②第二份 PDF 的作者页署名为 Nathan Beckmann, CMU（来源为 CSD faculty meeting Fall ‘21 的报告，被 15-418/618 复用为补充讲义）。
  • 已公开的姊妹课程（Stanford CS149）讲义抽取文本：已公开。 cs149_supp/accelerators.txt（Stanford CS149 Fall 2025 Lecture 10 “Hardware Specialization”，71 页）。它是本讲内容最完整的公开文本来源，覆盖了能效约束、专用化的量级、指令流开销分解、DSP/FPGA/ASIC/DSP、H100 SM 结构、张量核、TMA、脉动阵列、数据流架构与 kernel fusion。
  • 本笔记的事实基础：extracted/20_heterogeneity.txt（44 页全文抽取）+ extracted/20_specialization_csd.txt（53 页全文抽取）+ cs149_supp/accelerators.txt（71 页全文抽取）。讲义中没有给出可抽取数值的图表（如 Chung et al. MICRO 2010 的能效-面积曲线、Hill & Marty 的加速比曲线、Hameed et al. 的能耗堆叠柱状图）本笔记只沿用其定性结论与坐标轴含义；本笔记中所有带具体数字的算例都显式标注为”按讲义给定参数推导“，不冒充讲义原文数据。
  • 未发布 / 需登录：Fall 2026 日程表中 Lecture 18 一行的 slides/video 链接被 HTML 注释隐藏（注释原文为 “slides/video from a previous offering; uncomment when posted for Fall 2026”），因此本讲的录像（Panopto/YouTube）属未发布；Ed 讨论区、Autolab、Canvas 均需登录；历史学期位于 /afs/cs/academic/class/15418-*/public/ 之下的讲义（如 Performance Analysis/Profiling、Transactional Memory、AI in System Design 等）需要 CMU 登录，属未公开。
  • 课程语境：Fall 2026 授课教师为 Brian Railing 与 Dimitrios Skarlatos；课程由 Kayvon Fatahalian 创建。前一讲为 Lecture 17 “Transactional Memory”（Oct 5，Fall 2026 尚未发布讲义），后一讲为 Lecture 19 “Virtual Memory”（Oct 9，对应公开 PDF 13_virtualmemory.pdf）。本日日程另标注 “Assignment 3 due, Assignment 4 out”。

2. 核心概念与硬件/软件架构图解

2.1 起点：从 44× 到”该买哪种机器”

定义与目的：本讲以一个朴素但尖锐的对比开场。讲义指出，在 Assignment 1 里我们观察到：一个优化良好的、计算密集（compute-bound）的并行实现，比同一处理器上用 gcc -O3 编译的单线程 C 代码快约 44 倍。这个数字说明两件事：①并行化本身的收益是巨大的；②收益完全取决于程序的可并行比例。于是就有了本讲的第一个思想实验：

你要买一台新计算机系统。 处理器 A：4 个核，每核顺序性能为 P。 处理器 B：16 个核，每核顺序性能为 P/2。 系统其它部件完全相同。你选哪个？

直观解释（”它是什么？”）：把它想成搬砖。

• 处理器 A = 4 个壮汉，每人一次能抱 2 摞砖；
• 处理器 B = 16 个普通人，每人一次抱 1 摞砖。

总”人力”（芯片晶体管预算）完全一样。如果活儿能被平均切成 16 份（比如”把 1600 摞砖全搬到 100 米外”），16 个人赢；但如果这活儿里有一段只有一个人能干（比如”先由一个人决定砖该怎么码”），那 A 就用它的壮汉把这段快速干完，再去加入搬砖大军——而 B 只能让一个普通人在那儿慢慢决定，其余 15 人干等。这就是异构/非对称设计的全部动机：宁可牺牲一部分”并行吞吐资源”，也要保住”串行段的执行速度”。

数学形式（Hill & Marty 模型，讲义给出的资源受限 Amdahl 定律）：设

• f = 程序中可并行的比例；
• n = 总处理资源（例如芯片上的晶体管总量）；
• r = 分配给每个核的资源量，于是有 n/r 个核，每个核的顺序性能为 perf(r)。

讲义举例 n = 16：处理器 A 取 r_A = 4（4 个大核），处理器 B 取 r_B = 1（16 个小核），并把 perf(r) 建模为 √r（面积/资源翻倍，性能只提升 √2——这与后面 2.2 节的技术缩放模型完全一致），归一化 perf(1) = 1。

• 对称多核（symmetric）：全部核一样大 Speedup_sym(f, n, r) = 1 / ( (1−f)/perf(r) + f / ( (n/r)·perf(r) ) )
• 非对称多核（asymmetric）：1 个”胖核”（资源 r）+ (n−r) 个”瘦核”（资源 1） Speedup_asym(f, n, r) = 1 / ( (1−f)/perf(r) + f / ( perf(r) + n − r ) )

教学要点：非对称形式里，串行部分由胖核（perf(r)）承担，并行部分由”胖核 + 所有瘦核”一起承担。Beckmann 的讲义用 x = BigCore, T = PSmall + √x 表达同一个想法：大核在局部是”低效”的（每单位面积产出的吞吐低），但在全局是”高效”的（它抬高了整个程序的加速比上限）。补充讲义的原话是：”Big CPUs are locally inefficient, but globally efficient.”

表 2.1：Amdahl 定律的三种形态

2.2 技术缩放的三条曲线与暗硅（dark silicon）

定义与目的：要理解”为什么必须专用化”，必须先看清通用处理器性能增长为什么慢下来了。补充讲义（Beckmann）给出了一条非常干净的三层建模链条：

• 乐观模型：性能 = 晶体管数 × 频率 = N·f。因为 N' = 2N（摩尔定律：约每 18 个月晶体管翻倍）且”更小的晶体管也更快”f' = √2 f，所以每代可得 2√2 ≈ 2.83×。
• 悲观模型：如果完全无法利用新增晶体管（唯一处理器是单核且已经跑满），性能 = f' = √2 f，每代只有 1.41×。
• 现实模型：性能 ≈ √N · f（讲义注明”略微乐观”），每代 2×。直觉是”核数翻倍，但调度/串行/一致性开销吃掉一半”。

功耗的物理：Power = N × 单个晶体管的功耗，而单管功耗 = C·V²·f，其中电容 C ≈ 晶体管宽度、电压 V ≈ 晶体管长度、频率 f ≈ 1/晶体管长度。一步摩尔定律之后：N'=2N, C'=C/√2, V'=V/√2, f'=√2 f（面积 ×1/2，线度 ×1/√2）。于是：

Dennard 缩放 (1974–2003):
  Power' = N'·C'·V'²·f' = 2N · (C/√2) · (V/√2)² · (√2 f) = N·C·V²·f = Power
  ⇒ 功率恒定！即使晶体管多了一倍、翻转频率高了 √2 倍。
  ⇒ 多出来的晶体管可以"免费"全用上。

Dennard 缩放终结 (~2003，电压受能带隙等物理限制不能再降):
  Power' = N'·C'·V'²·f' = 2N · (C/√2) · V² · (√2 f) = 2 × Power   ← 功耗翻倍！
  但功率必须近似恒定 (散热器有限) ⇒ Power'/Power = 1
  ⇒ N'·f' = √2 · N·f
  ⇒ 可用的晶体管比例 = √2/2 ≈ 70.7%，剩下的 ~30% 每代都"点不亮" = 暗硅 (dark silicon)

摩尔定律没死，但另外两条缩放死了：讲义的原话是 “Moore’s Law isn’t dead — at least not in Taiwan & Korea. But! 成本缩放（cost scaling）停止了，Dennard（功耗）缩放停止了。” 这正是”暗硅 ⇒ 专用化”的逻辑起点（补充讲义引用了 Esmaeilzadeh et al. ISCA’11 “Dark silicon and the end of multicore scaling” 与 Venkatesh et al. ASPLOS’10 “Conservation cores”）。

直观解释（”它是什么？”）：把芯片想成一栋散热能力固定的公寓楼。

• Dennard 时代 = 每个房间的电器都逐年更省电，所以你可以往楼里不断加房间、加电器，总电费不变。
• Dennard 之后 = 电器不再变省电了，但楼里房间数每代翻倍。总电费有硬上限（散热器和电池），所以你必须永久性地关掉约 30% 的房间。那么问题就变成：这 70% 的房间，是全部做成”什么都能干但都不快”的通用房间，还是拿出一些做成”只干一件事但快 100 倍”的专用车间？

图解 2-1：移动端芯片的”功耗—时间”四道约束墙

讲义（heterogeneity, slide 24）给出的移动端处理经验法则是：一个任务运行得越久，它被允许使用的功耗就越低。原因是四道依次收紧的约束：

 芯片允许的功耗
   ^
   |  ① 供电上限 (electrical limit)  ─────────────────────────────────────
   |     ┌──────────────┐
   |     │              │
   |  ② │              │  结温上限 Tj ────────────────────────────────────
   |     │              │  (芯片超过 Tj 就不可靠; 高功耗只能短时间跑)
   |     │              │
   |  ③ │              │  机壳温度上限 ──────────────────────────────────
   |     │              │  (芯片本身温度还行, 但热已经传到外壳, 用户握不住了)
   |     │              │
   |  ④ │              │  电池寿命目标 ──────────────────────────────────
   |     │              │  (芯片和外壳都凉, 但为了续航必须压低平均功耗)
   |─────┴──────────────┴────────────────────────────────────────────────> 时间 t
         |←── 任务持续时间 (几毫秒 → 几秒 → 几分钟)→|
   讲义给出的电池容量参照: iPhone 6 ≈ 7 Wh; 9.7" iPad Pro ≈ 28 Wh; 15" MacBook Pro ≈ 99 Wh

这条曲线的工程含义：短任务可以”爆发式”用高频大核（比如按下快门那一瞬间的图像处理），长任务必须”细水长流”（比如持续视频播放、always-on 语音唤醒）。所以一台机器同时需要大核（吃爆发）和小核（吃长尾），也需要固定功能单元（把长任务的能耗压到 1/100）。

图解 2-2：通用处理器为什么”低效”——一条指令的完整生命周期，以及 H.264 编码的能耗去向

补充讲义（CS149）用一张”执行一条指令要经过哪些步骤”的清单回答了”通用处理器到底低效在哪”：

              一条标量 FP 指令在通用处理器中的完整生命周期
 ┌──────────────────────────────────────────────────────────────────────┐
 │ ① 取指: 读 I-cache → ② 译码 → ③ 译成 uops / 查 uop cache             │
 │ ④ 检查数据依赖与流水线冒险 (记分牌 / 寄存器重命名)                    │
 │ ⑤ 选择可用的执行资源 (调度 / 发射)                                    │
 │ ⑥ 用译码出的操作数去控制寄存器堆 (RF) SRAM 的读端口                   │
 │ ⑦ 把数据从寄存器堆搬到执行单元                                        │
 │ ⑧ 执行算术运算                      ← 只有这一步是"有用功"            │
 │ ⑨ 把结果搬回寄存器堆 → ⑩ 控制写回端口                                │
 │ ⑪ 地址翻译 (TLB / 页表) 配合 D-cache 访问 → ⑫ 实际的访存             │
 └──────────────────────────────────────────────────────────────────────┘
   每条指令都要把 ①–⑦、⑨–⑫ 的代价重付一遍。
   SIMD 的作用是让 ⑧ 变宽, 从而"摊薄"①–⑦、⑨–⑫ 的每元素开销;
   VLIW / 复杂指令 (DP4、4x4 MMA) 的作用是让一条指令里的 ⑧ 更多。

H.264 视频编码的实测能耗分解（Hameed et al. ISCA 2010，讲义以定性堆叠柱状图给出）显示：即使已经用 SIMD 指令实现了编码器，功能单元（FU）仍然只占能耗的一小片。

  H.264 视频编码的能耗分解 (Hameed et al. ISCA 2010)
   ┌────┐ ┌────┐ ┌────┐ ┌────┐   IF  = 取指 + I-cache
   │ IF │ │ IF │ │ IF │ │ IF │   Pip = 流水线级间寄存器
   ├────┤ ├────┤ ├────┤ ├────┤   Ctrl= 杂项流水控制
   │Pip │ │Pip │ │Pip │ │Pip │   RF  = 寄存器堆读写
   ├────┤ ├────┤ ├────┤ ├────┤   D-$ = 数据 cache
   │Ctrl│ │Ctrl│ │Ctrl│ │Ctrl│   FU  = 功能单元 ← 唯一算"有用功"的部分
   ├────┤ ├────┤ ├────┤ ├────┤
   │ RF │ │ RF │ │ RF │ │ RF │
   ├────┤ ├────┤ ├────┤ ├────┤
   │D-$ │ │D-$ │ │D-$ │ │D-$ │
   ├────┤ ├────┤ ├────┤ ├────┤
   │ FU │ │ FU │ │ FU │ │ FU │  ← 很小的一片
   └────┘ └────┘ └────┘ └────┘
    整数    亚像素   帧内预测  算术
   运动估计 运动估计 DCT/量化   编码
  ⇒ 通用处理器的"开销/有用功"比例极差 ⇒ Horowitz (ISSCC'14): CPU 只有 ≈1% 的能量花在有用计算上
                                                 ⇒ 这意味着 ~100× 的专用化机会

专用化的收益量级（讲义反复引用的三条经验法则，compared to high-quality C code on CPU）：

讲义用 “choosing the right tool for the job”（Credit: Pat Hanrahan 的这套分类图）总结：能效越高者越难编程，而”让专用硬件更好编程”本身是一个活跃的研究方向。

2.3 异构的形态：从 SoC 到超级计算机

定义与目的：异构并行处理（heterogeneous parallel processing）指在一台机器里混合多种”优化目标不同”的计算资源：延迟优化的顺序核（latency-optimized）、吞吐优化的并行核（throughput-optimized）、以及领域专用的固定功能单元（domain-specialized fixed-function）。讲义的原话是：

“Idea: the most efficient processor is a heterogeneous mixture of resources — use the most efficient tool for the job.”

直观解释（”它是什么？”）：现代 SoC 就像一个厨房。CPU 大核是主厨（什么菜都能做，但一次只做一道，单项速度最快）；CPU 小核是帮厨（切菜慢一点但省电，可以很多个一起干）；GPU 是流水线配菜工（同一道工序重复几千遍最快）；Neural Engine 是专用压面机（只会压面，但一秒压一万份，耗电极低）；视频编解码器是专用切片机。真正的设计难题不是”要不要买这些机器”，而是”这台机器上到底该配几台压面机、几个配菜工”——配少了，压面机成为瓶颈，几百个配菜工一起等它（讲义 slide 39 的”陷阱”）；配多了，面积被占用，通用处理能力下降。

讲义给出的真实系统实例（全部来自讲义原文）：

GPU 本身就是一个异构多核处理器：讲义画出 GPU 的内部结构——大量重复的 SIMD Exec + Cache 切片（这才是 Assignment 2 里 CUDA 程序用到的通用计算资源），外加一批图形专用的固定功能单元：Texture（贴图/滤波）、Clip/Cull、Rasterize、Tessellate、Zbuffer/Blend，以及最上层的 Scheduler / Work Distributor。讲义借此解释”什么是固定功能硬件加速的图形任务”：

• Rasterization（光栅化）：确定一个三角形覆盖了哪些像素；
• Texture mapping（纹理映射）：对图像做变形/滤波，把细节贴到表面上；
• Geometric tessellation（几何细分）：由粗糙几何计算出精细几何。

补充讲义（Beckmann）特别强调一个容易被误解的点：在”CPU + GPU”这个组合里，GPU 并不是这个故事里的 “accelerator”，它是一个可编程的 manycore（真正准确的说法是”介于二者之间”）。把 GPU 当”加速器”会让人误以为它只能做一件固定的事，从而忽略它作为通用吞吐平台的调度难题。

2.4 加速器设计原理：把能量全部花在有用功上

定义与目的：加速器（accelerator）的设计目标可以一句话概括——Spend all energy on useful work。做法是：

1. Replicate a minimal, custom circuit many times（把最小的定制电路复制很多份）：最大化并行、消除开销、回到”乐观缩放”的区间（面积与算力同比例增长，不被控制逻辑稀释）；
2. 数据值直接在运算单元之间路由，取消寄存器堆（Data values routed directly from one operation to another — no register file!）；
3. 配一大块片上存储把数值放在近处（Big on-chip memory / “scratchpad”）。

直观解释（”它是什么？”）：通用处理器像一个万能工位：工具放在远处的工具柜（寄存器堆/内存），每做一步都要”走过去拿工具、走回来用、再走过去放回”。加速器像一条固定工序的装配线：每台机器只做一件极小的事，工件直接从前一台机器滑到下一台，旁边放一大箱零件（scratchpad）随手可取。因为工序固定，不需要”看图纸”（取指译码），也不需要”每步都回工具柜”。

这就是”分工”的能效来源：讲义引用 Dally（NVIDIA/Stanford, SC’15）的结论——一旦把 CPU 的低效去掉，数据搬运就变成能耗的绝对主导。

图解 2-3：通用处理器 vs 专用加速器的结构对比

   通用处理器 (CPU / GPU core)                     专用加速器 (ASIC, 如 TPU/Eyeriss)
   每一步都经过寄存器堆 + 指令流                   数值由运算单元直接流向下一个运算单元

     指令流 (IF / ID / 调度)                          (无指令流: 无取指、译码、发射开销)
          │                                                    │
          v                                                    v
    ┌───────────────┐                                ┌───────────────────────┐
    │ 寄存器堆 RF   │<── 读/写端口 ──>                │ Scratchpad (片上 SRAM)│
    │ (面积大、能耗高)│    每个操作数都要搬进搬出       │ 大容量、靠得近         │
    └───────┬───────┘                                └────┬──────────┬───────┘
            │                                             │          │
    ┌───────v──┐ ┌──────────┐ ┌──────────┐          ┌─────v───┐  ┌───v─────┐
    │ ALU/FPU  │ │ ALU/FPU  │ │ ALU/FPU  │          │  PE 0   │─>│  PE 1   │─> ...
    │ (复用的) │ │ (复用的) │ │ (复用的) │          │ (MAC)   │  │ (MAC)   │
    └──────────┘ └──────────┘ └──────────┘          └─────────┘  └─────────┘
     少量"宽"的复用单元                                大量"窄"的复制单元 (replicate!)
     面积/能耗: 控制 >> 计算                            面积/能耗: 存储 >> 计算

   Dally (CACM'20) 的两句判据:
     "Achieving high speedups and gains in efficiency from specialized hardware
      usually requires modifying the underlying algorithm..."
     "When logic is free, memory dominates. Because the area and power of most
      accelerators are memory dominated, a reasonable first estimate of these costs
      can be made by considering only the memory..."
   ⇒ 加速器的第一性原理不是"算得快", 而是"用最少的搬运做同样多的计算"。

因此，加速器设计本质上是一种算法设计：要把算法改写成一个”数据局部性极高、中间结果尽量不落片外”的形式。讲义给出的两个正面例子：

• Eyeriss（Chen, Krishna, Emer, Sze, ISCA’16）：开创性的 CNN 加速器，特点是大容量片上存储 + 为 CNN 定制的片上网络拓扑。
• 脉动阵列（systolic array）：Google TPU 的核心，把”权重驻留 + 数据波前流动”做成了电路（见 2.5）。

2.5 脉动阵列（systolic array）：数据驱动的波前

定义与目的：脉动阵列是一组规则排列的处理单元（PE），每个 PE 只与邻居通信，数据像心脏搏动（systole）一样一拍一格地流过阵列。它的目的是把”矩阵乘”这种规则计算的全局访存替换成局部邻居通信，同时让权重”驻留”在 PE 内。TPU v1 用它做 y = Wx 与卷积。

直观解释（”它是什么？”）：把它想成一排排接力赛。

• 通用 SIMD 的做法像”一个教练（指令流）同时指挥 32 个人去仓库取货、各自算、再送回仓库“——每次都要过仓库（寄存器堆/内存）。
• 脉动阵列像”一条传菜流水线：食材（x）从流水线一端递进来，每经过一个工位就被该工位专属的调料（权重 w）处理一次，处理完就顺手传给下一个工位**“。传的是”手递手”（邻居连线），不用回仓库；结果（累加器 acc）就留在自己工位上。

图解 2-4：y = Wx 的 N×N 脉动阵列波前

   注入 x →  ┌───────┐   ┌───────┐   ┌───────┐   ┌───────┐
             │ PE 0  │──>│ PE 1  │──>│ PE 2  │──>│ PE 3  │
             │w00..  │   │w10..  │   │w20..  │   │w30..  │
             │w03    │   │w13    │   │w23    │   │w33    │   权重由顶部 Weights FIFO
             ├───────┤   ├───────┤   ├───────┤   ├───────┤   每拍向各 PE 推入一个,
             │ acc 0 │   │ acc 1 │   │ acc 2 │   │ acc 3 │   之后"驻留"在 PE 内
             └───┬───┘   └───┬───┘   └───┬───┘   └───┬───┘
                 │           │           │           │
                y0          y1          y2          y3
    每拍每个 PE 把 x 传给右邻 (局部连线), 不回寄存器堆、不回内存

   时间-空间图 (哪一拍哪个 PE 在做哪一次乘加):
     拍 t:      0      1      2      3      4      5      6
     PE 0:    w00    w01    w02    w03     -      -      -
     PE 1:     -     w10    w11    w12    w13     -      -
     PE 2:     -      -     w20    w21    w22    w23     -
     PE 3:     -      -      -     w30    w31    w32    w33
              └─────────── 填充 (fill) ───────────┘
                                                    ↑ y3 在 t = 6 就绪 ⇒ Span = 2N−1 = 7 拍
     稳态吞吐 = N 次 MAC / 拍 (每 PE 每拍 1 次 MAC)
     平均 PE 利用率 = N² MAC / (N 个 PE × (2N−1) 拍) = N/(2N−1) → 1/2

性能特征（延迟、带宽、吞吐量）：

• 吞吐量：稳态下每个 PE 每拍做一次 MAC，阵列吞吐 = N MAC/拍；TPU 用 256×256 的乘加阵列。
• 延迟 / Span：单个 y = Wx 的关键路径是 2N−1 拍（填充 N 拍 + 排空 N−1 拍）。
• 带宽：全局/片外搬运量从朴素实现的 O(N²) 降到 N² + N（权重只装载一次、x 只注入一次）；取而代之的是 (N−1)(2N−1) 次邻居间传递（N 个 PE 排成一线，每拍有 N−1 条链路上的传递，共 2N−1 拍；N=8 时是 105 次）——这看起来比搬运次数更多，但都是片上 1 mm 内的局部连线，能耗比 DRAM 低 1–2 个数量级（见 4.3 节 pJ 数值）。稳态下每次 MAC 只伴随约 (N−1)/N ≈ 1 次邻居传递。
• 效率：填充/排空期让单次运算的平均利用率约 50%；连续处理多个输入向量可以把排空期填满，利用率趋近 100%。

表 2.2：SIMD 与脉动阵列的对比（CS149 讲义原表）

2.6 软件执行模型：从 CUDA 到”分块张量 + 异步流水 + 算子融合”

定义与目的：硬件给了异构资源，软件必须提供能表达”数据搬运”的抽象。CS149 讲义把”理想 AI 模型加速器”该具备的特性列成了一张清单，并逐条给出为什么：

表 2.3：理想加速器的特性清单，以及 NVIDIA GPU 的完成度

分块张量为什么是”最大 TFLOPS”的钥匙：一条 16×16×16 的 tensor core 指令一次完成 4096 次乘加。CS149 讲义给出”可编程性开销”的估算——同样是完成一批乘加运算，半精度 FMA 形式的指令流开销约为 2000%，半精度 DP4（vec4 点积）降到 500%，半精度 4×4 MMA 只有 27%。核心原则是：

Key principle: amortize cost of instruction stream processing across many operations of a single complex instruction.（用一条复杂指令里的众多运算来摊薄指令流处理成本。）

这也解释了硬件数值格式的演进：BF16（S=1, E=8, M=7，与 FP32 同范围但精度更低）、BF8 E4M3（范围 0–448）、BF8 E5M2（范围 0–57344）——格式越来越窄，是因为单位面积/单位能量的吞吐才是目标（slide credit: Bill Dally）。

CUDA 的三层层级结构（CS149 给出的 H100 对照表）：

   CUDA 层级           计算层级              存储层级
   ──────────────────────────────────────────────────────────────────────
   Grid         ←→     GPU            ←→   80 GB HBM / 50 MB L2
   Cluster      ←→     CPC            ←→   每 SM 256 KB shared memory
   Thread Block ←→     SM             ←→   每 SM 256 KB shared memory
   Threads      ←→     SIMD Lanes     ←→   每线程 1 KB 寄存器,
                                          每 SM 分区 64 KB
   · Thread Block Cluster 是至多 16 个 thread block 的集合
   · 每个 thread block 保证跑在**独立的一个 SM** 上, 且**同时**执行
     (这是"计算单元到计算单元通信"的硬件基础)

图解 2-5：H100 的 Streaming Multiprocessor（SM）内部结构

  NVIDIA H100 SM  (讲义数据: 整芯片 144 个 SM; 每 SM 含 4 个 sub-core)
 ┌────────────────────────────────────────────────────────────────────────────┐
 │  Shared Memory / L1  (256 KB, 可配置)      ┌─────────────────────────────┐ │
 │  ┌────────────────────────────────────┐    │ Tensor Memory Accelerator   │ │
 │  │  片上暂存: 相当于"加速器的 scratchpad"│    │ (TMA): 单线程发起一整块张量 │ │
 │  └────────────────────────────────────┘    │ 的异步 global→shared 搬运,  │ │
 │                                             │ 由 copy descriptor 描述区域,│ │
 │  ┌────── Sub-core 0 ──────┐ ┌──── Sub-core 1 ────┐ │ 完成后用 barrier 通知 │ │
 │  │ Warp Selector          │ │ Warp Selector      │ └─────────────────────────────┘ │
 │  │ Fetch/Decode (1 warp/拍)│ │ Fetch/Decode       │                               │
 │  │ ┌────────────────────┐ │ │ ┌────────────────┐ │   每个 sub-core:              │
 │  │ │ 32-wide FP32 SIMD  │ │ │ │ 32-wide FP32   │ │   · FP32: 每拍 32 次 MUL-ADD │
 │  │ │ 32-wide INT32 SIMD │ │ │ │ 32-wide INT32  │ │   · INT32: 每 1 拍 16 次     │
 │  │ │ 16-wide FP64 SIMD  │ │ │ │ 16-wide FP64   │ │   · FP64: 每 2 拍 16 次      │
 │  │ │ LSU (load/store)   │ │ │ │ LSU            │ │   · 寄存器 64 KB / sub-core  │
 │  │ │ Tensor Core        │ │ │ │ Tensor Core    │ │   · 4 个 sub-core 共 256 KB  │
 │  │ │  16x16x16          │ │ │ │  16x16x16      │ │     寄存器, 由 ≤64 个 warp 分 │
 │  │ │  [fp16 × fp16→fp32]│ │ │ │  [fp16×fp16→fp32]│ │                          │
 │  │ └────────────────────┘ │ │ └────────────────┘ │                               │
 │  └────────────────────────┘ └────────────────────┘                               │
 └────────────────────────────────────────────────────────────────────────────┘
   同一张图上的两套算力 (A100/GA100 的讲义数字):
     SIMD 路径:   108 SM × 64 fp32 ALU × 2 flop × 1.4 GHz  = 19.5 TFLOP/s (fp32)
     Tensor Core: 108 SM × 4 个 tensor core                 = 312 TFLOP/s (fp16 输入, fp32 累加)
   ⇒ 94% 的峰值算力在专用单元里; H100 更进一步: tensor core 989 TFLOP/s (fp16),
     SIMD 只有 134 TFLOP/s (fp16) / 67 TFLOP/s (fp32) ⇒ "All the TFLOPS are in the
     Tensor Cores" (讲义给出的各代占比: 89% / 50% / 94% / 96% / 98%)。

软件侧的”异步流水”执行模型：为了不让加速器等数据，执行必须异步（non-blocking）——后一次搬运/计算在前一次完成之前就启动：

   主机(大核) 时间轴:  [准备 A ][准备 B ][准备 C ] ...     ← 串行/难并行部分留在 CPU
                          │        │        │
                          v        v        v
   加速器队列:      ┌────────────────────────────────────────────┐
                    │ LD0 │ ST0 │ AO0 │                          │   LD = 载入输入
                    │     │ LD1 │ ST1 │ AO1 │                    │   AO = 加速器上的算子
                    │     │     │ LD2 │ ST2 │ AO2 │              │   ST = 写回结果
                    └────────────────────────────────────────────┘
      关键: 后一次 LD/ST/AO 在前一次完成之前就已启动。
      被隐藏掉的是 DRAM 访问延迟与 PCIe/NVLink 传输时间 —— 不是算力。
      (CUDA 里的对应物: cp.async / TMA 异步搬运 + mbarrier + mma_async)

   数据流图 (AI 模型本身就是一张 dataflow graph):
     权重 ─┐
           ├──> [GEMM 1] ──> [Pool] ──> [GEMM 2] ──> [SoftMax] ──> [Sum] ──> 输出
     样本 ─┘
           每两个算子之间都要把中间结果写回片外、再读回来 ⇒ "GEMM 计算本身便宜,
           数据搬运才贵(硅面积、瓦特、纳秒)" ⇒ 于是要用 **算子融合 (fusion)**:
           FlashAttention 把 QK^T → Mask → Softmax → Dropout → ×V 融进同一条
           tile 级流水 (MetaPipeline), 中间结果根本不落片外。

图解 2-6：粗粒度可重构阵列（CGRA）——”把数据流图钉在硅上”

讲义（Beckmann）最后给出”通用加速”的一条路线：coarse-grained reconfigurable arrays (CGRAs)，代表系统是 Plasticine（Prabhakar et al., ISCA’17），后商业化为 SambaNova（讲义注明 2021 年 4 月估值 $5.1B；作者 Kunle Olukotun, Stanford）。它复用了 TRIPS、Garp、Piperench、Dyser 等一脉相承的老思想。

   粗粒度可重构阵列 (CGRA, 例: Plasticine)
   ┌──────────────────────────────────────────────────────────────────┐
   │   开关网络 (Switch) = 芯片上的"数据流布线"                        │
   │                                                                  │
   │   ┌────────┐   ┌────────┐         ┌────────┐   ┌────────┐        │
   │   │  PMU   │   │  PMU   │         │  PMU   │   │  PMU   │        │
   │   │ Pattern│   │ Pattern│         │ Pattern│   │ Pattern│        │
   │   │ Memory │   │ Memory │         │ Memory │   │ Memory │        │
   │   │ Unit   │   │ Unit   │         │ Unit   │   │ Unit   │        │
   │   └───┬────┘   └───┬────┘         └───┬────┘   └───┬────┘        │
   │       │            │                  │            │             │
   │   ┌───v────┐   ┌───v────┐         ┌───v────┐   ┌───v────┐        │
   │   │  PCU   │   │  PCU   │         │  PCU   │   │  PCU   │        │
   │   │ Pattern│   │ Pattern│         │ Pattern│   │ Pattern│        │
   │   │ Compute│   │ Compute│         │ Compute│   │ Compute│        │
   │   │ Unit   │   │ Unit   │         │ Unit   │   │ Unit   │        │
   │   │ (ALU 阵│   │ (ALU 阵│         │ (ALU 阵│   │ (ALU 阵│        │
   │   │  + op1 │   │  + op2 │         │  +pred │   │  +tid  │        │
   │   │  token │   │  token │         │  token │   │  token │        │
   │   │  store)│   │  store)│         │  store)│   │  store)│        │
   │   └────────┘   └────────┘         └────────┘   └────────┘        │
   └──────────────────────────────────────────────────────────────────┘
   处理器 = 复制出来的 PE 阵列 (复制是高效扩展的必要条件)
   编译器把指令映射到 PE 上, 并在整个执行期间**留在那里** ⇒ 直接编码数据流图
   ⇒ 没有取指/译码 (无指令流), 极端异步 (没有顺序指令执行), 只有算子间点对点搬运
   代价: 重新映射 = 重新编译整张图; 通用性的来源是"可重构", 而不是"可换指令"

2.7 异构系统的两大难题：配比与映射

定义与目的：讲义把异构的困难拆成”硬件设计者的难题”和”软件开发者的难题”两半。

硬件设计者的难题（配比问题）——讲义原话整理：

• 什么是正确的资源混合比（mixture of resources）才能同时满足性能、成本与能耗目标？
• 吞吐型资源太少 ⇒ 并行负载的峰值性能/效率低（本该把资源拿去做更多吞吐核）；
• 顺序处理资源太少 ⇒ 被 Amdahl 定律咬住；
• 该给某个特定功能（例如视频）分配多少芯片面积？（这些面积是从通用处理能力里扣出来的）；
• 如何把”有哪些单元”告诉软件？（抽象是什么？编译器和操作系统怎么办？）；
• 推论：芯片设计时必须极其准确地理解（未来的）工作负载 —— 系统无法随使用习惯变化、新算法出现而自适应；
• 趋势是保守地”超额配置”固定功能单元，从而削弱它们的优势。

软件开发者面临的难题（映射问题）：如何把程序映射到一堆异构资源上？

• “Pick the right tool for the job“：算法必须分解成若干部分，每一部分都恰好匹配机器上某一类计算单元；
• 异构系统的调度问题更复杂（要考虑数据传输、单元间依赖、谁的队列更长）；
• 可用的资源混合比会反过来决定算法的选择（如果机器上有一台 1000× 的 DNN ASIC，你可能就把算法改成 DNN 友好的形式）；
• 可移植性与维护是噩梦（”we’ll revisit this next class”）。

2.8 能耗：把”少搬数据”当成第一设计准则

定义与目的：讲义把减少能耗总结为两条思路：① 用专用化处理单元，② 少搬数据。

讲义给出的”ballpark”能耗数值（这是本讲最重要的量化工具表）：

表 2.4：数据搬运的能耗代价

直接推论（讲义原文）：

• 以 10 GB/s 从内存读数据 ⇒ 约 1.6 W 功耗；
• 而整个移动 GPU 的功耗预算只有约 1 W（别忘了手机还要跑 CPU、显示、射频等）；
• iPhone 6 电池只有 ~7 Wh（对比 MacBook Pro 笔记本的 99 Wh 电池）。

结论（讲义原话的意译）：“Exploiting locality matters!!!” 并且——“Suggests that recomputing values, rather than storing and reloading them, is a better answer when optimizing code for energy efficiency!”（为能效优化时，重新计算往往优于”存起来再读回来”。）

讲义总结的三条能效趋势：

1. Compute less!（少算）：计算要耗能，因此并行算法如果比串行版本做了更多的工作，即使跑得更快，也未必值得。但性能仍然重要，因为处理器只要开着就在烧”静态功耗”。
2. Specialize compute units（专用化计算单元）：异构处理器（CPU 类核 + 吞吐优化核）；固定功能单元（音频、”运动传感器处理”、视频编解码、图像处理/计算机视觉）；专用指令（不断扩张的 AVX 向量指令集、AES-NI 加密指令）；可编程软逻辑（FPGA）。
3. Reduce bandwidth requirements（降低带宽需求）：利用局部性（重构算法以尽量在片上复用数据）；激进地使用压缩（多做一点计算来压缩数据再写入内存，讲义预计会出现专门降低通用压缩/解压开销的固定功能硬件）。

3. 代码示例与性能分析

3.1 示例一：异构任务划分——把”可并行段”和”必须串行段”分给不同的单元

这个示例把 2.1 节的 Amdahl 思想写成一个可运行的程序：阶段 A 是完全数据并行、计算密集的（应该放到”多而小”的吞吐型单元上）；阶段 B 是一条串行递推（应该放到”少而快”的大核上）。程序实测两个阶段的时间占比，并用 Amdahl 定律预测不同核数下的加速比上限。

// 编译 (release):  g++ -O3 -fopenmp -march=native -std=c++17 het_pipeline.cpp -o het_pipeline
// 运行:            OMP_NUM_THREADS=16 ./het_pipeline
//
// 说明: -O3 打开向量化与循环优化; -march=native 让编译器用本机可用的 SIMD
//       指令集(SSE/AVX/AVX-512), 这一步对"阶段 A"的性能影响极大;
//       -fopenmp 提供 OpenMP 运行时(omp_get_wtime / omp_get_max_threads)。

#include <omp.h>
#include <cstdio>
#include <cstdlib>
#include <cmath>
#include <vector>

static const int H = 1024;   // 图像高
static const int W = 1024;   // 图像宽
static const int K = 32;     // 每个像素的迭代次数 (模拟计算密集程度)

// ---------------------------------------------------------------
// 阶段 A: 高度数据并行 (行与行之间完全独立)
// 映射建议: 放到"多而小"的吞吐型单元 (GPU / 小核 / 宽 SIMD 单元) 上。
// ---------------------------------------------------------------
static void stage_compute_parallel(const std::vector<float>& in,
                                   std::vector<float>& out)
{
    #pragma omp parallel for schedule(static)
    for (int y = 0; y < H; ++y) {
        for (int x = 0; x < W; ++x) {
            float v = in[(size_t)y * W + x];
            for (int k = 0; k < K; ++k) {
                // 带 sqrt 与 fabs 的非线性迭代: 延迟高、吞吐受限、无分支
                v = 0.999f * v + 0.001f * sqrtf(fabsf(v) + 1.0f) + 1e-6f;
            }
            out[(size_t)y * W + x] = v;
        }
    }
}

// ---------------------------------------------------------------
// 阶段 B: 顺序依赖强 (第 y 行的阈值依赖第 y-1 行的结果)
// 映射建议: 放到"少而快"的延迟优化核 (大核) 上 —— 大核在局部低效,
//           但在全局高效: 它把这段无法并行的拖尾时间压下去。
// ---------------------------------------------------------------
static float stage_adaptive_threshold(const std::vector<float>& img,
                                      std::vector<float>& mask)
{
    float thr = 0.5f;
    for (int y = 0; y < H; ++y) {
        double row_sum = 0.0;
        for (int x = 0; x < W; ++x) {
            float v = img[(size_t)y * W + x];
            mask[(size_t)y * W + x] = (v > thr) ? 1.0f : 0.0f;
            row_sum += v;
        }
        const float mean = (float)(row_sum / W);
        thr = 0.9f * thr + 0.1f * mean;   // 串行递推: 无法并行, 也无法向量化
    }
    return thr;
}

int main()
{
    std::vector<float> in((size_t)H * W), out((size_t)H * W), mask((size_t)H * W);
    for (size_t i = 0; i < in.size(); ++i)
        in[i] = 0.5f + 0.5f * sinf((float)i * 1e-4f);

    const double t0 = omp_get_wtime();
    stage_compute_parallel(in, out);
    const double t1 = omp_get_wtime();
    const float thr = stage_adaptive_threshold(out, mask);
    const double t2 = omp_get_wtime();

    const double ta = t1 - t0, tb = t2 - t1, tt = ta + tb;
    printf("threads            = %d\n", omp_get_max_threads());
    printf("stage A (parallel) = %8.3f ms  (%5.1f%% of total)\n",
           ta * 1e3, 100.0 * ta / tt);
    printf("stage B (serial)   = %8.3f ms  (%5.1f%% of total)\n",
           tb * 1e3, 100.0 * tb / tt);
    printf("total              = %8.3f ms   (final threshold = %.4f)\n",
           tt * 1e3, thr);

    // Amdahl: 用实测的串行时间占比 s 预测"再加核"的收益上限
    const double s = tb / tt;
    printf("\ns = serial fraction = %.3f\n", s);
    for (int p = 1; p <= 64; p *= 2)
        printf("  P=%2d cores -> speedup ceiling = %6.2fx   (asym. ceiling = %6.2fx)\n",
               p, 1.0 / (s + (1.0 - s) / p), 1.0 / s);
    return 0;
}

【代码做什么？】

1. main 先构造一张 1024×1024 的浮点”图像”（用 sinf 生成有空间变化的输入，避免常量折叠）。
2. stage_compute_parallel：对每一行做 32 次非线性迭代。每个像素的迭代次数固定、无分支、无跨像素依赖，因此这一阶段是”教科书式的数据并行”。
3. stage_adaptive_threshold：逐行计算已处理图像的均值，并用 thr ← 0.9·thr + 0.1·mean 更新阈值。第 y 行用的是第 y−1 行算出的阈值，这是一条真正的串行递推（类似 IIR 滤波），无法用并行前缀（scan）简单拆开，也无法向量化（阈值是循环携带依赖）。
4. 用 omp_get_wtime() 分别给两个阶段计时，打印各阶段的时间占比 s，再用 Amdahl 公式 1/(s + (1−s)/P) 打印不同核数下的加速比上限，同时打印 1/s 这一”核数无穷大时的天花板”。
5. 编译时必须用 -O3：不开优化时阶段 A 的 sqrtf/循环开销会淹没真实结构，测出的”串行比例”会失真。

【并行机制与性能解说】

• 硬件上怎么并行：#pragma omp parallel for schedule(static) 让 OpenMP 运行时在进入循环时唤醒一个线程池（线程数 = OMP_NUM_THREADS，默认为物理核数），把 y ∈ [0, H) 的迭代空间静态切块（默认连续切块），每个线程拿一段连续行。因为每次迭代访问独立的内存地址（行内连续，行间相隔 W 个 float），不存在写共享与伪共享（false sharing）：不同线程写的缓存行不重叠（W = 1024 个 float = 4096 B，正好是 64 B 缓存行的整数倍）。若把 W 改成非对齐的奇数，相邻行会共享缓存行，写共享缓存行会触发缓存行在两个核之间反复迁移（ping-pong），性能可能掉数倍——这是把”行切分”改成”列切分”或任意 1D 切分时的经典陷阱。

• 阶段 B 为什么留在单线程：它是一条长度为 H·W 的依赖链，每个元素的写入都依赖前一个元素的中间状态（row_sum 与 thr）。如果强行并行（例如每个线程负责若干行、各自维护局部阈值），结果与串行语义不等价，属于”为了并行而改了算法”。若确实要并行，正确做法是改成两遍（two-pass）：第一遍并行算所有行的均值，第二遍并行算 thr 的串行递推——把 thr 的递推单独抽成一条长度为 H 的小串行链（Span 从 H·W 降到 H）。

• Work / Span / 并行度（按 H = 1024, W = 1024, K = 32 计算）：

  量	阶段 A	阶段 B	合计
  Work（总工作量）	H·W·K = 3.355 × 10⁷ 次迭代	H·W = 1.049 × 10⁶ 次逐元素更新	3.460 × 10⁷
  Span（关键路径）	W·K = 3.277 × 10⁴（一行内串行）	H·W = 1.049 × 10⁶（完全串行，并行度 1）	1.081 × 10⁶
  并行度 = Work/Span	H = 1024	1	3.460×10⁷ / 1.081×10⁶ ≈ 32.0

  解读：整程序的理论并行度只有约 32，尽管阶段 A 单独看有 1024 的并行度。这正是异构机器存在的理由——把一个”局部并行度 1024、局部串行度 1”的程序，拆给两种不同的单元。可扩展性上限由阶段 B 决定：即便并行部分无限快，整体加速比也不会超过 Work/Span = 32×。

• 数值化的性能预测（按假设参数推导）：设阶段 A 的单次迭代在开了 SIMD 的核上平均耗 10 个周期、CPU 主频 3.0 GHz、16 个线程；阶段 B 的每次逐元素更新耗 4 个周期、单线程。
  • 单线程总时间：3.355e7 × 10 / 3.0e9 + 1.049e6 × 4 / 3.0e9 = 111.8 ms + 1.40 ms = 113.2 ms；
  • 16 线程总时间：111.8/16 + 1.40 = 6.99 + 1.40 = 8.39 ms；
  • 实测加速比 ≈ 113.2/8.39 = 13.5×，而 Amdahl 上限（s = 1.40/113.2 = 1.24%）是 1/(0.0124 + 0.9876/16) = 13.6×（吻合），核数→∞ 的天花板是 1/0.0124 ≈ 81×。
  • 结论：在这个例子里 s = 1.2% 尚可接受；但如果阶段 B 的常数因子再大 10 倍（s ≈ 11%），16 核加速比就掉到 1/(0.11 + 0.89/16) = 6.0×——“给串行段配一个大核”（提高 perf(r)）比”再堆 8 个小核”有效得多，这就是 2.1 节非对称模型的直接体现。

3.2 示例二：CUDA 分块矩阵乘（tiled GEMM）——”当逻辑免费时，存储说了算”

这个示例是 2.6 节”分块张量 + 片上复用”思想的最小可运行版本：用共享内存（shared memory，即 GPU 上的 scratchpad）把全局内存流量降低 TILE 倍。

// 编译 (release):  nvcc -O3 -arch=sm_80 --use_fast_math tiled_gemm.cu -o tiled_gemm
//                  (也可以在 nvcc 前加 -Xptxas -O3; 不开 -O3 时内层循环不会充分展开)
// 运行:            ./tiled_gemm
//
// 说明: TILE 决定片上复用度。TILE=16 时一个 block 是 16x16=256 个线程,
//       每个线程算 C 的一个元素, 累加器留在寄存器里。

#include <cstdio>
#include <cstdlib>
#include <cmath>
#include <cuda_runtime.h>

#define TILE 16

#define CHECK(call)                                                          \
    do {                                                                     \
        cudaError_t e_ = (call);                                             \
        if (e_ != cudaSuccess) {                                             \
            fprintf(stderr, "CUDA error %s:%d: %s\n", __FILE__, __LINE__,    \
                    cudaGetErrorString(e_));                                 \
            exit(1);                                                         \
        }                                                                    \
    } while (0)

// C = A * B, 三个都是 N x N 的行主序方阵
__global__ void tiled_gemm(const float* __restrict__ A,
                           const float* __restrict__ B,
                           float* __restrict__ C, int N)
{
    __shared__ float As[TILE][TILE];
    __shared__ float Bs[TILE][TILE];

    const int tx = threadIdx.x;          // 0..TILE-1
    const int ty = threadIdx.y;          // 0..TILE-1
    const int row = blockIdx.y * TILE + ty;
    const int col = blockIdx.x * TILE + tx;

    float acc = 0.0f;                    // 累加器常驻寄存器, 不落片外

    for (int t = 0; t < N / TILE; ++t) {
        // 1) 协作式搬运: 每个线程搬 1 个 A 元素和 1 个 B 元素到片上 scratchpad
        As[ty][tx] = A[(size_t)row * N + (t * TILE + tx)];
        Bs[ty][tx] = B[((size_t)(t * TILE + ty)) * N + col];
        __syncthreads();                 // 2) 保证 tile 装载完成

        // 3) 在片上 scratchpad 里做 TILE 次乘加 (无全局访存)
        #pragma unroll
        for (int k = 0; k < TILE; ++k)
            acc += As[ty][k] * Bs[k][tx];

        __syncthreads();                 // 4) 防止下一轮覆盖本轮的 tile
    }
    C[(size_t)row * N + col] = acc;
}

int main()
{
    const int N = 1024;
    const size_t bytes = (size_t)N * N * sizeof(float);

    float *hA = (float*)malloc(bytes), *hB = (float*)malloc(bytes),
          *hC = (float*)malloc(bytes);
    for (size_t i = 0; i < (size_t)N * N; ++i) { hA[i] = 1.0f; hB[i] = 1.0f; }

    float *dA, *dB, *dC;
    CHECK(cudaMalloc(&dA, bytes));
    CHECK(cudaMalloc(&dB, bytes));
    CHECK(cudaMalloc(&dC, bytes));
    CHECK(cudaMemcpy(dA, hA, bytes, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(dB, hB, bytes, cudaMemcpyHostToDevice));

    const dim3 block(TILE, TILE);
    const dim3 grid(N / TILE, N / TILE);

    tiled_gemm<<<grid, block>>>(dA, dB, dC, N);   // 预热: 排除首次启动开销
    CHECK(cudaDeviceSynchronize());

    cudaEvent_t e0, e1;
    CHECK(cudaEventCreate(&e0));
    CHECK(cudaEventCreate(&e1));
    CHECK(cudaEventRecord(e0));
    const int ITER = 20;
    for (int it = 0; it < ITER; ++it)
        tiled_gemm<<<grid, block>>>(dA, dB, dC, N);
    CHECK(cudaEventRecord(e1));
    CHECK(cudaDeviceSynchronize());

    float ms = 0.0f;
    CHECK(cudaEventElapsedTime(&ms, e0, e1));
    ms /= ITER;

    const double flops = 2.0 * (double)N * N * N;          // N^3 次 MAC = 2N^3 FLOP
    const double bytes_moved = 4.0 * 2.0 * (double)N * N * N / TILE;  // 见下文推导
    printf("N=%d  TILE=%d  kernel = %.3f ms  ->  %.1f GFLOP/s\n",
           N, TILE, ms, flops / (ms * 1e-3) / 1e9);
    printf("  DRAM traffic (analytic) = %.2f GB, achieved BW = %.0f GB/s\n",
           bytes_moved / 1e9, bytes_moved / (ms * 1e-3) / 1e9);
    printf("  arithmetic intensity    = %.2f FLOP/byte\n", flops / bytes_moved);

    CHECK(cudaMemcpy(hC, dC, bytes, cudaMemcpyDeviceToHost));
    printf("C[0]=%f (期望 %d)\n", hC[0], N);   // A 全 1, B 全 1 => 每行和 = N
    CHECK(cudaFree(dA)); CHECK(cudaFree(dB)); CHECK(cudaFree(dC));
    free(hA); free(hB); free(hC);
    return 0;
}

【代码做什么？】

1. 主机端分配三个 N×N（N = 1024）的方阵；A、B 全填 1，因此正确答案是 C[i][j] = N，便于校验。
2. 用 dim3 block(TILE, TILE)（256 个线程）和 dim3 grid(N/TILE, N/TILE)（4096 个 block）启动核函数：一个 thread block 负责 C 的 TILE×TILE 一小块。
3. 核函数主循环沿 k 维以 TILE 为步长推进。每轮：
   • 每个线程从全局内存各取一个 A 元素和一个 B 元素写入共享内存 As/Bs（协作式搬运）；
   • __syncthreads() 保证整块 tile 都装好再使用（否则会有线程读到未写入的共享内存）；
   • 在片上做 TILE 次乘加（#pragma unroll 让 16 次迭代全展开，编译器可把它们排成流水，用多个 FMA 发射槽掩盖 FMA 的延迟）；
   • 第二次 __syncthreads() 防止快线程进入下一轮、在本轮数据被用完前就覆盖 As/Bs（这是”共享内存 + 双同步”模式的核心正确性要求）。
4. 计时用 cudaEvent（GPU 侧时间戳，比主机 clock() 准确），先跑一次预热，再跑 20 次取平均。
5. 程序同时打印解析求出的 DRAM 流量与据此算出的算术强度，便于与 Roofline 模型（见 4.4 节）对照。

【并行机制与性能解说】

• 硬件上怎么并行：grid = (64, 64) 个 block 被分发到 GPU 的多个 SM 上；每个 block 内的 256 个线程被组织成 8 个 warp（每 32 线程一个 warp），由 SM 的 warp scheduler 每拍发射一个 warp。对 As[ty][k] * Bs[k][tx] 这条语句，同一 warp 内 32 个线程的 ty/tx 取值使它们访问不同的共享内存地址；由于共享内存被划分为 32 个 bank，若同一 warp 的 32 个线程落在 32 个不同 bank 上就是无冲突（conflict-free）的一拍访问；若地址对齐不当（例如 TILE 是 32 的倍数且按列访问），多个线程会落到同一 bank，产生 bank conflict，把一次访问串行化成多拍。这是 tiled GEMM 里最常见的性能陷阱之一（工程上的标准修法是给 As 的行宽加一个 padding，例如 __shared__ float As[TILE][TILE + 1]）。

• 数据类型与 __restrict__：__restrict__ 告诉编译器三个指针互不别名，使编译器可以对全局访存做更激进的调度；若去掉它，编译器可能因为”C 可能别名 A“而不敢重排访存。

• Work / Span / 并行度：

  量	值	说明
  Work	N³ = 1.074 × 10⁹ 次 MAC（= 2N³ = 2.147 × 10⁹ FLOP）	与算法无关的总工作量
  Span（关键路径）	N = 1024 次 MAC	单个输出元素的累加链长度（每轮 tile 1 次 + 线程内串行累加）；这是任何分块方案都改不掉的算法下界
  并行度 = Work/Span	N² = 1.049 × 10⁶	恰好等于”同时活跃的线程总数”（N/TILE × N/TILE × TILE²）

  解读：并行度百万量级，远超任何 GPU 的物理线程容量，所以瓶颈绝不在并行度。真正的瓶颈是每个活跃线程要为 1 次 MAC 运多少字节——也就是算术强度。这解释了为什么加速器设计的第一性原理是”减少数据搬运”而不是”增加并行度”。

• 工作量与数据搬运的定量关系（按 N = 1024, TILE = 16, fp32 推导）：网格有 (N/TILE)² 个 block；每个 block 需要读一整行 panel A[TILE × N] 与一整列 panel B[N × TILE]，即 2·TILE·N 个元素。总 DRAM 流量（每轮 tile 都从全局内存重新取，理想化为无 L2 命中）：

  元素数 = (N/TILE)² × 2·TILE·N = 2N³/TILE
         = 2 × 1.074e9 / 16 = 1.342e8 个元素
  字节数 = 1.342e8 × 4 B = 0.537 GB
  算术强度 AI = 2N³ FLOP / (4 × 2N³/TILE) B = TILE/4 = 4.0 FLOP/byte
  

  这是低估（没有算 L2 命中，真实 DRAM 流量会更小），但结论方向正确：AI = TILE/4 只与分块边长成正比。A100 的 fp32 峰值 19.5 TFLOP/s、HBM 带宽 ~1555 GB/s，其 ridge point = 19.5e12 / 1555e9 ≈ 12.5 FLOP/byte。TILE = 16 时 AI = 4.0 < 12.5，核函数是访存受限的：理论最短时间 = 0.537 GB / 1555 GB/s = 0.345 ms（用 N=1024 的规模），而计算只需 2.147e9/19.5e12 = 0.110 ms——访存时间是计算的 3.1 倍。若把有效分块做到 TILE = 64（需要”寄存器分块”：每个线程算 4×4 = 16 个输出，用 256 个线程覆盖 64×64 的块），AI 升到 16 > 12.5，核函数转为计算受限，此时才可能接近峰值。若再上 tensor core（fp16 输入、fp32 累加，A100 上 312 TFLOP/s），ridge point 升至 312e12/1555e9 = 200 FLOP/byte，需要多层复用 + 算子融合（TMA 异步搬运 + 大 tile + 把相邻算子融进同一条流水）才可能喂饱——这正是讲义”异步访存 + 计算单元间通信 + 融合”这张清单的由来。

3.3 示例三：脉动阵列模拟——把”全局访存”换成”邻居通信”

这个示例用纯 C 实现 2.5 节的 y = Wx 脉动阵列，并统计两种实现的访存次数，从而把”专用化换来的是什么”变成可数的量。

/* 编译: gcc -O3 -std=c11 systolic.c -o systolic
 * 运行: ./systolic
 *
 * 说明: -O3 让内层小循环完全展开; 这个程序的目的是"数访存次数",
 *       所以刻意不做任何缓存优化, 让每次装载都被显式计数。 */

#include <stdio.h>
#include <stdlib.h>
#include <math.h>

#define N 8
#define TMAX (2 * N - 1)      /* 波前深度: 填充 N 拍 + 排空 N-1 拍 */

/* 全局计数器: 从"权重/输入存储"(视作片外或全局共享存储)装载一个数 */
static long g_loads = 0;
/* 全局计数器: PE 与 PE 之间的局部搬运次数 */
static long g_pe_hop = 0;
/* 全局计数器: 直接作用于累加器的访存次数(读改写) */
static long g_acc_rmw = 0;

/* ---------------- 朴素实现: 每个输出都独立地从存储取操作数 ---------------- */
static void naive(const float W[N][N], const float x[N], float y[N])
{
    for (int i = 0; i < N; ++i) {
        float acc = 0.0f;
        for (int j = 0; j < N; ++j) {
            g_loads  += 2;    /* 读 W[i][j] + 读 x[j]      */
            g_acc_rmw += 1;   /* 累加器读改写               */
            acc += W[i][j] * x[j];
        }
        y[i] = acc;
    }
}

/* ---------------- 脉动阵列: 权重驻留 + x 向右流动 ---------------- */
static void systolic(const float W[N][N], const float x[N], float y[N])
{
    float acc[N];
    float xr[N + 1];          /* 波前寄存器: xr[i] 是 PE i 本拍看到的 x 值 */
    float wf[N][N];           /* 权重 FIFO 的快照(每个权重只搬运一次)      */

    for (int i = 0; i <= N; ++i) xr[i] = 0.0f;
    for (int i = 0; i < N; ++i) {
        acc[i] = 0.0f;
        for (int j = 0; j < N; ++j) wf[i][j] = W[i][j];
    }
    g_loads += (long)N * N;   /* 权重在整个计算期间只从存储搬运一次 */

    for (int t = 0; t < TMAX; ++t) {
        if (t < N) { xr[0] = x[t]; g_loads += 1; }  /* 每个 x 只装载一次 */
        else       { xr[0] = 0.0f; }                /* 排空期注入 0     */

        for (int i = 0; i < N; ++i) {
            const int j = t - i;            /* PE i 在 t 拍处理 W[i][j] * x[j] */
            if (j >= 0 && j < N) {
                acc[i] += wf[i][j] * xr[i]; /* 权重在 PE 内; x 来自左邻寄存器 */
                g_acc_rmw += 1;
            }
            if (i < N - 1) g_pe_hop += 1;   /* 本拍把值传给右邻 */
        }

        /* 波前向右推进一格(倒序拷贝, 保证读到的是本拍的旧值) */
        for (int i = N - 1; i >= 0; --i) xr[i + 1] = xr[i];
    }

    for (int i = 0; i < N; ++i) y[i] = acc[i];
}

int main(void)
{
    float W[N][N], x[N], y_naive[N], y_sys[N];
    for (int i = 0; i < N; ++i) {
        x[i] = 1.0f + 0.1f * (float)i;
        for (int j = 0; j < N; ++j) W[i][j] = (float)((i + 1) * (j + 1)) * 0.01f;
    }

    naive(W, x, y_naive);
    const long loads_naive = g_loads, hop_naive = g_pe_hop, rmw_naive = g_acc_rmw;

    g_loads = g_pe_hop = g_acc_rmw = 0;
    systolic(W, x, y_sys);
    const long loads_sys = g_loads, hop_sys = g_pe_hop, rmw_sys = g_acc_rmw;

    double maxerr = 0.0;
    for (int i = 0; i < N; ++i)
        maxerr = fmax(maxerr, fabs((double)y_naive[i] - (double)y_sys[i]));

    printf("N = %d, 波前深度 Span = 2N-1 = %d 拍, 总 MAC 数 = %d\n\n",
           N, TMAX, N * N);
    printf("%-28s %10s %10s\n", "计数器", "朴素实现", "脉动阵列");
    printf("%-28s %10ld %10ld\n", "从存储装载次数", loads_naive, loads_sys);
    printf("%-28s %10ld %10ld\n", "累加器读改写次数", rmw_naive, rmw_sys);
    printf("%-28s %10ld %10ld\n", "PE 间局部搬运次数", hop_naive, hop_sys);
    printf("%-28s %10ld %10ld\n", "存储系统访问总计", loads_naive + rmw_naive,
           loads_sys + rmw_sys);
    printf("\n最大误差 = %.3e (两实现结果一致)\n", maxerr);
    return 0;
}

【代码做什么？】

1. naive 按 y[i] = Σ_j W[i][j]·x[j] 实现，每做一次 MAC 计 2 次”装载”（读权重、读 x）与 1 次”累加器读改写”。
2. systolic 用一个长度为 N+1 的 xr[] 数组模拟”波前寄存器”：xr[i] 表示 PE i 本拍看到的 x 值。每拍先在最左端注入 x[t]（排空期注入 0），再让所有 PE 做一次 MAC（PE i 用 W[i][t−i]），最后把 xr[] 倒序向右搬一格。权重 wf[N][N] 只是在开始时装载一次（计入 N²），此后驻留在 PE 内。
3. 两次运行分别累加三个全局计数器：存储装载次数、累加器读改写次数、PE 间局部搬运次数，最后并排打印，并校验两种实现结果一致（最大误差应为 0）。

【并行机制与性能解说】

• 硬件上怎么并行：N 个 PE 空间上并排放置，每拍同时做一次 MAC。它不是”同一条指令作用在 32 条通道上”（SIMD），而是”N 个独立的小电路各自持有自己的权重、各自累加“，控制是分布式的（每个 PE 只需要”权重 FIFO 给我下一个 w”和”左邻给我下一个 x”两个事件）。这正是 2.5 节对比表中”数据驱动（wavefront）”、”局部通信”、”分布式控制”的硬件含义。

• Work / Span / 并行度（N = 8）：

  量	值
  Work	N² = 64 次 MAC
  Span（波前深度）	2N−1 = 15 拍
  并行度 = Work/Span	64/15 ≈ 4.27（峰值可用并行度是 N = 8 个 PE）
  稳态利用率	N² / (N·(2N−1)) = 64/(8×15) = 53%；N 越大越趋近 N/(2N−1) → 50%

  解读：Work/Span ≈ 4.27 小于 PE 数 8，说明填充/排空（fill/drain）就是效率上限，不是算术量不够。要让利用率接近 100%，必须连续喂入多个 x 向量（把前一次的排空期用后一次的填充期填满）——这正是真实加速器用流水线/多批次（batching）掩盖启动开销的一般手法。

• 访存行为的对比（N = 8，即程序实际输出的数字）：朴素实现 2N² = 128 次装载 + N² = 64 次累加器读改写 = 192 次存储系统访问；脉动阵列 N² + N = 72 次装载 + N² = 64 次本地累加 = 136 次，外加 (N−1)(2N−1) = 105 次 PE 间局部搬运。按表 2.4 的能耗数值估算单次 MAC 的代价：

  朴素 (操作数来自 1mm 外的片上 SRAM, 26 pJ / 64 bit):
      2 次装载 × 26 pJ + 1 次累加读改写 × 26 pJ + FP 乘加 20 pJ =  98 pJ / MAC
  朴素 (操作数实际来自 LPDDR, 1200 pJ / 64 bit):
      2 × 1200 pJ + 20 pJ                                        ≈ 2420 pJ / MAC
  脉动阵列 (权重驻留、x 来自左邻连线、acc 在本地寄存器):
      稳态: 每次 MAC 伴随 (N-1)/N ≈ 1 次邻居传递
            1 × 26 pJ (按片上 1mm 内 SRAM 的上界估) + 20 pJ       =  46 pJ / MAC
      含填充/排空: 实测 105 次传递 / 64 次 MAC = 1.64 次每 MAC
            1.64 × 26 pJ + 20 pJ                                  ≈  63 pJ / MAC
  ⇒ 同样 64 次 MAC: 朴素(命中 DRAM) ≈ 155 nJ; 脉动阵列 ≈ 4.0 nJ (含填充/排空),
    稳态下则只有 ≈ 2.9 nJ ⇒ 约 39× (稳态约 53×) 的能效差
  

  这正是讲义”CPU 只有约 1% 的能量花在有用计算上 ⇒ ~100× 机会”的微观版本：把”回寄存器堆/回内存取操作数”换成”邻居手递手”，能耗就能掉一到两个数量级。注意这里的代价是牺牲通用性——这个阵列只会算 y = Wx，遇到分支、不规则访存、稀疏结构就完全失效。

4. 性能模型与复杂度分析

4.1 模型一：技术缩放——为什么”每代只能点亮 70% 的晶体管”

假设与推导（补充讲义 Beckmann 给出的代数）：设一代之后 N' = 2N（摩尔定律）、C' = C/√2、V' = V/√2、f' = √2 f（几何缩放，面积减半、线度乘 1/√2）。

表 4.1：三种性能模型的数值对比（归一化初始代 = 1）

Dennard 缩放（1974–2003）：Power' = N'C'V'²f' = 2N·(C/√2)·(V/√2)²·(√2 f) = N·C·V²·f = Power ⇒ 功耗恒定，多出来的晶体管可以全部用上。

Dennard 终结后（~2003 起）：电压不能再按 1/√2 下降（能带隙等物理限制），于是

Power' = N'·C'·V'²·f' = 2N·(C/√2)·V²·(√2 f) = 2·Power        ← 功耗翻倍
但散热约束要求 Power' / Power = 1
⇒ N'·f' = √2 · N·f
⇒ 可用于"新增算力"的晶体管比例 = √2 / 2 = 70.71%
⇒ 每代约 29.3% 的新晶体管无法被点亮 = 暗硅 (dark silicon)

数值算例（暗硅的规模）：假设 Gen0 有 N = 10⁹ 个晶体管、频率 f = 1.0，则 Gen1 有 2 × 10⁹ 个晶体管，但功率预算只允许 N'f' = √2 × 1.0 = 1.414。若把频率保持不变（f' = f = 1.0），可用晶体管数只有 1.414 × 10⁹，即 2 × 10⁹ 中有 5.86 × 10⁸ 个（29.3%）必须长期关断。若选择”频率提高 √2 倍”（f' = 1.414），则只能同时点亮 1.0 × 10⁹ 个晶体管——新增的一整代晶体管（10⁹ 个）全部变成暗硅，性能只涨 1.41×。这就是”必须专用化”的定量压力：既然不能把所有晶体管都做成通用的，那就把一部分做成只在特定任务上开机、但能效高 100 倍的专用单元。

4.2 模型二：资源受限的 Amdahl 定律与”非对称为什么赢”

用 2.1 节的公式，取 n = 16、perf(r) = √r、perf(1) = 1。

表 4.2：加速比数值表（按公式计算，n = 16，perf(r) = √r）

（示例展开 f = 0.90, r = 4： 对称 = 1/(0.10/2 + 0.90/(4×2)) = 1/(0.05 + 0.1125) = 6.15； 非对称 = 1/(0.10/2 + 0.90/(2 + 16 − 4)) = 1/(0.05 + 0.064286) = 8.75。 f = 0.90, r = 2 非对称 = 1/(0.10/1.41421 + 0.90/(1.41421+14)) = 1/(0.070711+0.058388) = 7.75。 f = 0.99, r = 2 对称 = 1/(0.01/1.41421 + 0.99/(8×1.41421)) = 1/(0.007071+0.087500) = 10.57。）

从表里能读出四条结论（这几条就是本讲前一半的骨架）：

1. 对称设计没有一个”正确的核大小”：三种 f 下对称设计的最优 r 分别是 1 / 1 / 1，但曲线很陡——f = 0.99 时从 r = 1 的 13.91× 掉到 r = 8 的 5.60×，跌了一半多。讲义的原话是 “No good single choice for core size!”
2. 非对称设计在”串行段不可忽略”时明显更好：f = 0.90 时非对称最优 8.75×（r = 4）比对称最优 6.40× 高 37%；f = 0.95 时 10.77× 比 9.14× 高 18%；曲线也更平坦（r 从 2 到 8 都能保持在 8.4× 以上）。讲义原话：”Speedup is much better & less sensitive to core size!”
3. 但当 f → 1 时，非对称的优势会缩小到几乎消失：f = 0.99 时对称最优 13.91×、非对称最优 14.03×，只差 0.9%。原因是此时串行项（(1−f)/perf(r) = 0.0071）已经小于并行项（0.0642），”把资源花在胖核上”换来的是并行吞吐的净损失。这是一个重要的现实提醒：异构设计的收益与工作负载的串行比例强相关，所以讲义才反复强调”芯片设计时必须准确理解工作负载”，而”系统无法随使用习惯和新算法自适应”。
4. 两者都仍受并行段限制：r = 16（只剩一个大核）时无论对称还是非对称都掉到 4.00×——把资源全押在串行性能上同样会输。异构的正解是”一个胖核 + 尽可能多的小核”，而不是”全胖”或”全瘦”。

补充讲义还给出了 manycore 版本的对比：纯 manycore 的加速比是 √x / ((1−s) + s/P)（x = 核大小，P = n/x），并注明 “Speedup is surprisingly awful unless program is ≥99.9% parallel!“。这与 4.5 节加速器时代的 99.9% 门限是同一个数学。

4.3 模型三：能效的”算术强度”与能耗墙

把表 2.4 的 pJ 数值变成一个可计算的能量模型：

E = N_ops × e_op + N_bytes × e_byte
  e_op    = 20 pJ / FLOP        (浮点运算)
  e_byte  = 1200 pJ / 8 B = 150 pJ/B   (LPDDR)
  e_byte  =   26 pJ / 8 B = 3.25 pJ/B  (片上 SRAM, 1 mm 距离)

能量平衡点 (计算能耗 = 搬运能耗):
  AI* = e_byte / e_op = 150 / 20 = 7.5 ops/byte = 15 FLOP/byte   (LPDDR)
  AI* = 3.25 / 20 = 0.1625 ops/byte = 0.325 FLOP/byte            (片上 SRAM)
  ⇒ 片上数据的"能量价格"只有 DRAM 的 26/1200 ≈ 1/46

数值算例（按讲义给定的 pJ 数值推导）：一个 256 MiB 的 float32 数组（2.684 × 10⁸ 字节，即 6.711 × 10⁷ 个元素），做单次遍历（每元素 1 次 FMA），内存带宽 20 GB/s。

遍历字节数         = 256 MiB = 2.684 × 10⁸ B
64 位(8 B)读取次数  = 2.684e8 / 8 = 3.355 × 10⁷ 次
搬运能耗           = 3.355e7 × 1200 pJ = 4.026 × 10⁻² J = 40.3 mJ
计算能耗           = 6.711e7 × 20 pJ = 1.342 × 10⁻³ J = 1.34 mJ
搬运 : 计算 能耗比  = 40.3 / 1.34 ≈ 30 : 1
                  (等价地: 每 8 个元素搬一次的能耗 1200 pJ,
                   对比 8 次 FMA 的 160 pJ ⇒ 7.5 : 1;
                   此时算术强度只有 2 FLOP / 4 B = 0.5 FLOP/byte)
理论最短时间       = 2.684e8 B / 20 GB/s = 13.42 ms
平均功率           = 40.3 mJ / 13.42 ms ≈ 3.0 W   (> 整个移动 GPU 约 1 W 的预算)

（若按十进制把 256 MB 解释为 2.56 × 10⁸ 字节，则最短时间为 12.8 ms，平均功率约 3.1 W——结论不变。）

讲义给出的同源数字：以 10 GB/s 从内存读数据 ≈ 1.6 W（用本笔记参数推导：10e9/8 × 1200 pJ = 1.5 W，与讲义吻合）。结论：移动端一个”每字节只算一次”的 kernel，光是把数据搬进来就已经超过整块 GPU 的功耗预算——所以优化能效的第一手段是提高算术强度（分块、融合、压缩），而不是提高频率或加核。

4.4 模型四：Roofline——”分块到底要多大才够”

Roofline 模型把可达性能写成

Attainable(FLOP/s) = min( Peak, AI × BW )
AI = 算术强度 (FLOP/byte)
ridge point = Peak / BW      (超过这个强度才算"计算受限")

表 4.3：两台 GPU 的 ridge point（用讲义给出的峰值与公开的 HBM 带宽）

数值算例（承 3.2 节的分块 GEMM，N = 4096）：

Work = 2N³ = 2 × 6.87e10 = 1.374 × 10¹¹ FLOP

分块 T=16 (每线程 1 个输出, 累加器在寄存器):
  DRAM 元素 = 2N³/T = 2 × 6.87e10 / 16 = 8.59 × 10⁹ 个 float
  DRAM 字节 = 3.44 × 10¹⁰ B = 34.4 GB
  AI        = 1.374e11 / 3.44e10 = 4.0 FLOP/byte   (< 12.5 ⇒ 访存受限)
  访存时间  = 34.4 GB / 1555 GB/s = 22.1 ms
  计算时间  = 1.374e11 / 19.5e12  =  7.05 ms
  ⇒ 实际时间 ≈ 22.1 ms (若无重叠), 等效算力 = 1.374e11/22.1e-3 = 6.2 TFLOP/s = 峰值的 32%

分块 T=64 (寄存器分块: 每线程 4x4 = 16 个输出):
  DRAM 字节 = 34.4 GB / 4 = 8.59 GB
  AI        = 16.0 FLOP/byte  (> 12.5 ⇒ 计算受限)
  访存时间  = 5.52 ms  <  计算时间 7.05 ms
  ⇒ 可以逼近 19.5 TFLOP/s 的理论峰值, 只是需要足够的 occupancy 来隐藏延迟

若要喂饱 tensor core (ridge = 200 FLOP/byte):
  需要 AI ≥ 200, 即片上有效复用度 ≥ 200/(2 B/元素) = 100 个元素级复用
  ⇒ 单级 shared memory 分块做不到 (128×128 的 fp16 分块只有 AI = 64 FLOP/byte)
  ⇒ 必须叠加: 多级 tiling (shared → register/TMEM) + TMA 异步搬运 +
     算子融合 (把 Softmax、Mask、Dropout 融进同一条 tile 流水)

这一段量化分析就是讲义那句 “When logic is free, memory dominates” 的数字版本：A100/H100 里 94%–98% 的峰值算力都在 tensor core 上，而要把这些算力兑现，唯一的路是把数据搬运最少化——分块、异步、融合、片上暂存（scratchpad/TMEM），也就是本讲所有专用化手段的共同名字。CS149 讲义用 “Hardware Lottery”（Sara Hooker）描述这一正反馈循环：TPU 擅长稠密矩阵乘（OI ∝ n）→ 人们设计 Transformer 这种”矩阵乘友好”的模型 → Transformer 主导 → 硬件进一步为矩阵乘专用化。

4.5 模型五：加速器时代的 Amdahl 定律——为什么门限是 99.9%

假设：某加速器把覆盖到的那部分代码加速 S_acc 倍。整体加速比仍由 Amdahl 给出：Speedup = 1 / ((1 − f) + f/S_acc)。

表 4.4：1000× 加速器下的覆盖率敏感度（按公式计算）

这就是讲义里两次出现的同一句话：“To scale, ≥99.9% of the program must run on the accelerator (déjà vu…)” —— 与 manycore 时代 ≥99.9% parallel 的门限完全同构。

数值算例（多个加速器的边际收益，按公式计算）：基线时间 = 1，逐步往 SoC 上加加速器：

读表要点：

• 第一个加速器虽然最快（1000×），但因为剩余 5% 的拖尾，整体只拿到 19.6×——加速器的绝对倍数不等于整体收益；
• 第二个加速器只有 100×，边际收益 1.95×；第三个只有 10×，边际收益反而高达 7.1×，因为它消灭了当前的瓶颈项（0.025 / 0.0262 ≈ 95% 的剩余时间都在它身上）；
• 讲义因此提出两个尖锐的问题：“How many accelerators do we need?（收益迅速递减，但前期工程成本巨大）” 和 “Are accelerators a scaling strategy or a one-off boost?（加速器是可扩展的战略，还是一次性的提升？）”；
• 表里没有计入的加速器自身开销（启动延迟、主机↔设备搬运、同步）恰恰是现实中最常吃掉收益的部分（见第 6 节）。

4.6 模型六：编程性与专用化的权衡清单

表 4.5：加速器时代的开放问题（补充讲义原话整理）

5. 关键要点

1. 异构的第一性原理是”用最合适的工具干活”，而不是”让所有单元都忙”。 同构系统的最优策略是”让每个处理器一直忙”，异构系统的最优策略是”让每个任务落在最合适的单元上“。这要求算法能被分解成”高度数据并行段 / 难以并行段 / 宽 SIMD 友好段 / 分支发散段 / 可预测访存段 / 不可预测但缓存友好段”，进而映射到 CPU 大核、CPU 小核、GPU、DSP、固定功能单元上。这个分解工作被推给了程序员，这是本讲最重要的”代价”。

2. 能效已经取代峰值性能成为第一设计约束，而专用化是提高能效的唯一大幅手段。 Dennard 缩放终结 ⇒ 功耗恒定而晶体管翻倍 ⇒ 每代约 29.3% 的晶体管成为暗硅。经验量级：吞吐优化核（GPU）~10× perf/W、域专用加速器 ~20×、FPGA ~50×（学界仍有争议）、固定功能 ASIC ~100–1000× perf/W（ASIC 用 ~1/1000 的面积、~1/100 的能量达到一个 CPU 核的性能）。代价是：能效越高，越难编程。

3. 专用化的收益不是”算得更快”，而是”搬得更少”。 一旦去掉通用处理器的指令/寄存器堆开销（CPU 只有约 1% 的能量花在有用计算上），数据搬运就成为能耗主导：浮点运算 ~20 pJ，而读 64 位片外 DRAM 要 ~1200 pJ（片上 1 mm 处的 SRAM 只要 ~26 pJ）。因此加速器的结构特征——scratchpad 大容量片上存储、数据在运算单元之间直接路由、取消寄存器堆、权重驻留、脉动波前、算子融合、异步搬运（TMA/TMEM/mbarrier）——全都是同一条原则的不同实现：把复用做到极致，让片外流量最小化。

4. “核该多大”没有全局最优解，但”一个胖核 + 尽可能多的小核”通常最优。 在 perf(r) = √r、n = 16 的资源受限 Amdahl 模型下，对称设计的加速比对核大小敏感且最优值随 f 漂移（f = 0.95 时 16 个小核给 9.14×，4 个胖核只给 6.96×）；非对称设计（1 个 r = 4 的胖核 + 12 个小核）在 f = 0.95 时给到 10.77×，且曲线平坦。大核”局部低效、全局高效”——它抬高的是整个程序加速比的上限（Amdahl 的串行分母）。

5. 专用化的门限极陡：99.9%，而且在加速器时代会再次出现。 对 1000× 的加速器，覆盖 90% 只得到 9.9×，覆盖 99% 得到 91.0×，覆盖 99.9% 才到 500.3×。因此工程上有两条推论：①加速器必须做成”平台/流水线”而不是”孤立的函数替换”（否则非库代码立刻成为瓶颈）；②多个加速器的部署顺序应当按”当前瓶颈项”来排，而不是按”加速倍数最大的排前面”。

6. 常见陷阱与注意事项

• 陷阱一：以为”异构 = 更快”，忽略了配比错误会反过来拖垮整机。 讲义给出的经典反例（Molnar 2010）：假设工作负载中 10% 是光栅化，你给固定功能光栅化单元分配了 1% 的芯片面积，而实际需要的是 1.2%。于是光栅化成为瓶颈，占芯片 99% 的可编程处理器全部空等——讲义的原话是”99% of the chip runs at 80% efficiency!”（吞吐只剩 1/1.2 ≈ 83%，与”约 80%”一致）。因为固定功能单元一旦流片就无法加速（”Work balance must be anticipated at chip design time”），业界的趋势是保守地超额配置，而这会削弱专用化的优势。教训：异构系统的性能由最弱的那个专用单元决定，而不是由最强的那个决定。

• 陷阱二：把”像素级并行”写成了”列切分”，踩上伪共享（false sharing）。 在 3.1 节的图像并行例子里，如果按列分块（每个线程负责若干列），相邻线程会在同一缓存行上写不同的字：每次写都要把该行升级为独占，导致缓存行在两个核之间来回迁移（ping-pong），实测可能掉数倍性能。正确做法是按行（或按 ≥64 B 对齐的分块）切分，或给每线程数据加 padding 使其独占缓存行。同理，在 CUDA 里 __shared__ float As[TILE][TILE] 按列访问会出现 bank conflict（一个 warp 内多个线程落到同一 bank），标准修法是 padding 成 [TILE][TILE + 1]。

• 陷阱三：忘了共享内存/scratchpad 的”双屏障”约束。 3.2 节的 __syncthreads() 少了任何一个都会出错：少了装载后的那次同步，线程会读到别人还没写入的 tile；少了计算后的那次同步，跑得快的 warp 会在别人还没用完 tile 时就把它覆盖掉（数据竞争）。这是 GPU 上”用片上暂存换带宽”模式的必付正确性成本，也是程序员最容易因”去掉一次同步看着更快”而引入的 bug。

• 陷阱四：只用算术强度（或只用带宽）单指标评估，忽略另一侧的天花板。 Roofline 的两侧都要算：TILE = 16 的 GEMM 在 A100 上 AI = 4 FLOP/byte < 12.5 的 ridge point，因此再怎么优化指令调度都是在访存墙上撞；反过来，把 TILE 提到 64 之后 AI = 16，此时若 occupancy 不足（寄存器压力导致活跃 warp 太少），又无法隐藏访存延迟。“带宽受限”和”延迟受限”要分开诊断，不要一概归因于”内存不足”。

• 陷阱五：忽略”重新计算 vs 存下来再读”的能耗不对称，以及多算一点反而更快的现象。 讲义明确指出：为能效优化时，重新计算往往优于”存起来再读回来”；同时”计算更少”也是趋势之一——一个比串行版本做更多工作的并行算法，即使更快，也未必是好的选择（因为处理器开着就在烧静态功耗）。这与”并行算法常常用冗余计算换并行度”的直觉相反，需要在能效目标下重新审视（例：N-body 用更贵的算法换更少的通信，可能反而省能量）。

• 陷阱六：低估加速器的固定开销，让 Amdahl 的漂亮数字落空。 4.5 节算出”三个加速器 → 270×”，但这个模型没有计入：主机↔设备的数据搬运（PCIe/NVLink 带宽通常远低于 HBM）、kernel 启动延迟、队列同步、以及”非库代码”的拖尾。讲义因此把”如何把可用加速器发布给软件（OS 抽象）“和”非库代码会成为 Amdahl 瓶颈“列为两个未解决的开放问题。工程上的对策是把加速器组织成异步流水线（LD/ST/AO 重叠，见 2.6 节的图），让搬运时间与计算时间互相隐藏，而不是让它们串起来相加。

7. 思考题（带答案）

思考题 1：用资源受限的 Amdahl 定律决策”买哪种机器”

题目：设 n = 32 单位芯片资源，perf(r) = √r，程序的 f = 0.98。有三个候选设计： (a) 32 个小核（r = 1）； (b) 8 个中等核（r = 4，对称）； (c) 1 个胖核（r = 8）+ 24 个小核（r = 1，非对称）。 分别计算加速比，并解释为什么 (c) 最优。如果程序的 f 降到 0.90，结论会变吗？

【答案】

• (a) 对称 r = 1：perf(1) = 1，核数 = 32/1 = 32。 S = 1/((1−0.98)/1 + 0.98/(32×1)) = 1/(0.02 + 0.030625) = 1/0.050625 = 19.75×
• (b) 对称 r = 4：perf(4) = 2，核数 = 32/4 = 8。 S = 1/(0.02/2 + 0.98/(8×2)) = 1/(0.01 + 0.06125) = 1/0.07125 = 14.04×
• (c) 非对称 r = 8：perf(8) = √8 = 2.8284，胖核 1 个 + 瘦核 32 − 8 = 24 个。 S = 1/(0.02/2.8284 + 0.98/(2.8284 + 24)) = 1/(0.007071 + 0.036520) = 1/0.043591 = 22.94×

结论：(c) 最优（22.94×），比纯小核 (a) 好 16%，比对称胖核 (b) 好 63%。 原因：f = 0.98 意味着串行部分占 2%，这在 19.75× 的加速比下已经是主导项（0.02 对比并行项 0.030625 占了 39.5% 的剩余时间）。非对称设计用 perf(8) = 2.83 把串行项从 0.02 压到 0.007071（降 64.6%），同时只牺牲 8 份资源（并行项从 0.030625 升到 0.036520，升 19%）——用 19% 的并行吞吐损失换来 65% 的串行时间削减，净赚。这就是”大核局部低效、全局高效”。

若 f = 0.90：

• (a) 1/(0.10 + 0.90/32) = 1/(0.10+0.028125) = 7.80×
• (b) 1/(0.10/2 + 0.90/16) = 1/(0.05+0.05625) = 9.41×
• (c) 1/(0.10/2.8284 + 0.90/26.8284) = 1/(0.035355+0.033546) = 14.51×

结论仍然 (c) 最优，但已不是”押注大核”：此时 (a) 最差（7.80×）的原因恰恰是”串行项 0.10 相对并行项 0.028 太大”，而 (b) 之所以反超 (a)，是因为它用一个 perf(4) = 2 的核把串行项减半（0.10 → 0.05）而并行项损失不大。这说明最优的资源混合比强烈依赖 f——这正是讲义”芯片设计时必须极其准确地理解工作负载”以及”系统无法随使用习惯和新算法自适应”的量化根据。工程上唯一的缓解办法是让架构在宽范围内都”不太差”（(c) 在两种 f 下都是最好的，曲线平坦），而不是针对某一个 f 精确调优。

思考题 2：给一个真实的 kernel 做”专用化收益”预算

题目：某移动端应用在一颗通用 CPU 小核上耗时 100 ms，时间分布为：DNN 推理 80 ms、视频解码 15 ms、其余（音频处理 + 控制逻辑 + 数据搬运）5 ms。已知：移动 GPU 能效是 CPU 核的 10×（且推理能完美映射到数据并行），固定功能视频解码器能效是 CPU 的 200×，音频 DSP 是 CPU 的 20×。请计算：①全用 GPU 加速推理的整体加速比；②再加上视频解码器的整体加速比；③再加上音频 DSP 的整体加速比；④用表 2.4 的能耗数值估算”音频那 5 ms 里如果有 3 ms 是数据搬运，改成片上复用后能省多少能量”（假设搬运的是从 LPDDR 读 64 位、频率为每 64 位一次搬运、3 ms 内共搬运 200 MB）。

【答案】

一条必须遵守的记账规则：Amdahl 定律里的 f 必须始终以原始基线时间为 1 来定义。把”已经被加速过的时间”再拿去做分母，是这类题最常见的错误。下面每一步都用原始占比 0.80 / 0.15 / 0.03 / 0.02 代入（音频 3 ms 与控制逻辑 2 ms 需从原来的 5 ms 里拆出来）。

①只加速 DNN 推理（GPU，10×）：S = 1/(0.80/10 + 0.15 + 0.03 + 0.02) = 1/(0.08 + 0.20) = 1/0.28 = 3.57×，总时间 100/3.57 = 28.0 ms。

②再加视频解码器（200×）： S = 1/(0.80/10 + 0.15/200 + 0.03 + 0.02) = 1/(0.08 + 0.00075 + 0.05) = 1/0.13075 = 7.65×，总时间 13.08 ms。 边际收益 7.65/3.57 = 2.14×——因为此时真正的瓶颈已经变成视频解码（0.15 是剩余时间 0.20 中的 75%）。

③再加音频 DSP（20×）： S = 1/(0.80/10 + 0.15/200 + 0.03/20 + 0.02) = 1/(0.08 + 0.00075 + 0.0015 + 0.02) = 1/0.10225 = 9.78×，总时间 10.23 ms。 边际收益 9.78/7.65 = 1.28×。此时剩余时间已经全部落在推理后的 0.08与不可加速的 0.02 上，继续加加速器的空间已经很有限（上限 1/0.02 = 50×）。

④音频段数据搬运的能耗节省：200 MB = 2.0 × 10⁸ B，按 8 B 一次读数计，共 2.5 × 10⁷ 次 LPDDR 访问。

• 从 LPDDR 搬运（1200 pJ/次）：2.5e7 × 1200 pJ = 3.0 × 10⁻² J = 30.0 mJ；
• 若改成从片上 SRAM（26 pJ/次 64 位）复用：2.5e7 × 26 pJ = 6.5 × 10⁻⁴ J = 0.65 mJ；
• 节省 ≈ 29.35 mJ，降低约 46 倍（1200/26 ≈ 46.2）。
• 对照 iPhone 6 的 7 Wh 电池（7 × 3600 = 25200 J）：单看 29.35 mJ 微不足道，但如果这是”always-on”功能、每秒重复 10 次，那就是 0.294 W——接近整块移动 GPU 的功耗预算（约 1 W）的 30%。这解释了讲义为什么给”always-on 语音唤醒”配一个专用 ASIC：不是因为它快，而是因为只有把每次检测的能量压到微焦级，才能一直开着。

小结：①单独加速占比最大的那一块收益可观但受 Amdahl 限制（10× 的单元只换来 3.57×）；②加速非瓶颈部分是浪费——如果先加了 200× 的解码器再去加 GPU，几乎拿不到收益，因为瓶颈已经换了位置；③真正的收益来自”按顺序消灭当前瓶颈”，这与 4.5 节的边际收益表完全一致（那里的第三个单元只有 10× 却贡献了 7.1× 的边际收益）；④能效优化的杠杆不只是”换更快的单元”，更在于”减少数据搬运”——第 ④ 问显示把片外访问换成片上复用可以直接把这一段的能量砍掉约 46 倍，而这与加速单元的倍数无关。

思考题 3：什么时候不该用专用硬件？以及”多少个加速器才够”

题目：你的团队正在设计一颗面向”通用桌面应用”的 SoC，候选的加速单元有：DNN 加速器（100×，覆盖典型工作负载的 60%）、视频编解码器（100×，覆盖 10%）、图像处理 ISP（50×，覆盖 5%）、安全/加解密引擎（30×，覆盖 3%），其余 22% 是控制逻辑与不可加速的通用代码。 (a) 全部加上后，整体加速比是多少？ (b) 依次累加四个单元，各步的整体加速比与边际收益分别是多少？从中能得出什么部署顺序的建议？ (c) 如果这 22% 的”通用代码”里有 15% 其实是某个库函数之外的、自研的图算法，用现有 DSL 无法表达，你会怎么处理？请结合讲义的开放问题回答。 (d) 讲义说”tendency is to be conservative, and over-provision fixed-function components (diminishing their advantage)”——请用本讲的光栅化反例解释这句话的含义，并说明为什么”超额配置”在商业上是理性的。

【答案】

(a) 全部加上：基线 = 1.00，各段原始占比 0.60 / 0.10 / 0.05 / 0.03 / 0.22。 S = 1/(0.60/100 + 0.10/100 + 0.05/50 + 0.03/30 + 0.22) = 1/(0.006 + 0.001 + 0.001 + 0.001 + 0.22) = 1/0.229 = 4.37×

(b) 逐步累加：

建议：①收益迅速递减（2.46 → 1.32 → 1.19 → 1.13），第 2 个之后的边际收益都在 1.2× 左右，而每个加速器的”前期工程成本巨大”（讲义：”Large up-front engineering costs”）——所以加加速器不该按”谁加速倍数大”排序，而该按”谁覆盖的时间多”排序，并且要在边际收益跌破门槛时停手；②本例中最终被 0.22 的通用部分彻底锁死（哪怕前四项全部做到无穷快，上限也只有 1/0.22 = 4.55×，而当前已经拿到 4.37×，后两个单元合起来只贡献了 4.5% 的加速比余量）。这是”加速器是可扩展战略还是一次性提升”这个问题的量化答案：一旦通用部分成为约束，继续堆加速器就完全不经济。

(c) 自研图算法的处理：讲义明确指出 DSL 路线的两个问题——“Problem: Non-library code becomes Amdahl bottleneck” 与 “Problem: What if an important operation is missing?”。本例里那 15% 正是这两个问题同时发作：即便把所有加速器都加上，0.22 的通用部分也会把整体加速比死死压在 4.55× 以下。可行的对策（按本讲给出的方向）：

1. 改算法而不只是改硬件：Dally 的话——”从专用硬件取得高加速比与高能效通常需要修改底层算法“。把图算法重构成”能用框架懂的高层操作（如稀疏 GEMM、scan、segmented reduce）表达”的形式，把它拉进已有加速器的覆盖范围；这也是讲义强调 “the available mixture of resources can dictate choice of algorithm” 的含义。
2. 做领域专用而不是函数专用：为图计算单独配一个域专用单元（讲义在 SoC 卡通图里就画了一个 “Graph processor?”），把覆盖比例从”库函数”提升到”领域”。
3. 走通用加速路线（CGRA / 可重构数据流）：讲义把 CGRA 列为”general-purpose acceleration”的方向——无指令流 ⇒ 无取指/译码开销，极端异步，用”重新映射”代替”重新取指”。代价是一次映射等于一次编译，且每个算子会被”钉”在固定的硅上。
4. 接受它，并把它做成”不阻碍流水线”的部分：如果这 15% 必须留在 CPU 上，那就把它与加速器流水重叠（异步 LD/ST/AO），确保它不成为串行瓶颈；同时把它的数据搬运降到最少（3.3 节/陷阱五：重新计算优于存取）。

(d) “保守超额配置”的含义：光栅化反例说明固定功能单元一旦配少了就无法补救——你需要在设计时预判 10% 的工作负载需要 1.2% 的面积，但只给了 1%，结果占芯片 99% 的可编程处理器全部空等（讲义原话 “99% of the chip runs at 80% efficiency!”）。代价是乘性的：不是”光栅化慢 20%”，而是”整块芯片的有效吞吐掉 20%”，且没有任何软件手段能救（”System cannot adapt to changes in usage over time, new algorithms, etc.”）。所以从风险管理角度，给固定功能单元留出超额余量是理性的——即便这会浪费面积、削弱专用化本应带来的优势（这正是讲义所说的 “diminishing their advantage”）。这也解释了为什么现代 SoC 会选择“可编程 + 专用”的混合：把”可能变化的量”交给可编程单元（GPU/DSP/CGRA），只把”绝不会变的量”（视频编解码格式、加密原语、传感器前端）交给固定 ASIC，从而在面积效率与适应未来负载变化的能力之间取得平衡。