附录 A：Exam 1 复习笔记（Lecture 1–13）

目录 · ← l26

附录 A：Exam 1 复习笔记（Lecture 1–13）

依据官方公开讲义 lectures/revision_exam1.pdf（24 页，首页标注 Fall 2023/Fall 2024 版本，属讲义沿用）整理，并对照 lectures/01–13 的公开讲义原文补全。 考试形式：当堂闭卷，可用 一张 A4 双面手写纸，必须黑/蓝笔，无计算器/电子设备；题型以简答与选择+解释为主（解释的分值远高于选项）。

A.1 官方复习幻灯片覆盖清单（逐页）

A.2 核心公式一张表

A.3 高频考点一：ISPC / SPMD 与 SIMD

官方复习题（slide 2–3 原文）：给出 sinx 的 two versions，要求分析交错 vs 分块分配。

交错（interleaved）版本：

gang 有 programCount 个实例（如 AVX2 下为 8）
  实例 0 处理 idx = 0, 8, 16, 24, ...
  实例 1 处理 idx = 1, 9, 17, 25, ...
  ...
  → 每次对 idx 的访问是一段连续的 8 个 float
  → 生成 packed load / packed store（单条 vmovups 搬 32 字节）

分块（blocked）版本：

  实例 0 处理 idx = 0 .. N/8−1（连续 128 个元素）
  实例 1 处理 idx = N/8 .. 2N/8−1
  ...
  → 同一时刻 8 个实例访问的地址相距 N/8 个 float
  → 生成 gather / scatter（AVX2 的 vgatherdps 极慢，约比 packed load 慢一个数量级）

结论：交错分配在 SIMD 上通常更优，因为一个 gang 的并发访问是连续的，能够合并（coalesce）成整宽向量访存。这也解释了为什么 ISPC 的 foreach 默认采用交错分配。

必答要点：

• programCount：gang 中同时执行的实例数（uniform 值），由目标 ISA 的向量宽度决定（SSE 4、AVX2 8、AVX-512 16）。
• programIndex：当前实例在 gang 中的 id（varying 值）。
• uniform：类型修饰符，所有实例共享同一值，纯粹是优化提示，不影响正确性。
• gang 执行模型：所有实例执行同一条指令流（coherent execution），分支导致发散（divergence），硬件用掩码串行化各分支路径。

常考陷阱：uniform 不是正确性要求（写成 varying 也正确，只是慢）；programCount 与 programIndex 不能用于标量代码。

A.4 高频考点二：CUDA 线程层次与同步

官方复习题（slide 4）：gridDim、blockIdx、blockDim、threadIdx 各是什么？为什么没有 gridIdx 和 threadDim？

• 四个内建变量：
  • gridDim：整个 grid 的维度（block 数量）
  • blockIdx：当前 block 在 grid 中的索引
  • blockDim：一个 block 的维度（线程数量）
  • threadIdx：当前线程在 block 内的索引
• 全局线程 id：int i = blockIdx.x * blockDim.x + threadIdx.x;
• 为什么没有 gridIdx：一个 kernel 只由一个 grid 启动（”launch a grid of CUDA thread blocks”），因此 grid 层之上没有”第几个 grid”的概念；grid 是启动单位而不是层级中的一层。
• 为什么没有 threadDim：线程是层级的最底层，其下没有子层级，所以不需要”我的维度”这个概念；反之 block、grid 都有子元素，因此它们有 dim。

同步三件套（slide 6）：

1. __syncthreads()：block 内屏障，等待同一 block 的所有线程到达。不能跨 block。
2. 原子操作：atomicAdd/atomicCAS/atomicExch/...，可用于 global memory 与 shared memory。
3. kernel 返回时的隐式全局屏障：所有线程结束时 kernel 才算完成，主机侧 cudaDeviceSynchronize() 或拷贝回数据时等待。CUDA 不提供 kernel 内部的全局（跨 block）屏障——这正是”用 kernel 分解代替全局同步”这一设计模式的原因。

经典错误：依赖 block 间顺序（例如让 block 0 写、block 1 读，用标志位自旋等待）。两个 block 可能被调度到同一个 SM 并占满资源，导致”自旋的 block 等着永不被调度的 block”⇒ 死锁。正确做法是拆成两个 kernel。

A.5 高频考点三：共享地址空间求解器（SPMD）

官方复习题（slide 7）：给出二维网格求解器代码，问其中的错误/性能问题。逐项排查：

红黑着色解法：把格点按 (i+j) 的奇偶分为两组。同色格点的 5 点 stencil 只依赖异色格点，因此同色格点之间无依赖，可完全并行；一次 sweep 只需两个相位 + 一次屏障。

A.6 高频考点四：性能模型（每考必有）

必须能在 30 秒内手算的量：

例 1（计算峰值）：16 核 × 3.0 GHz × 8 宽 AVX2 × 2（FMA）= 768 GFLOP/s。 例 2（带宽上限）：某程序算术强度 4 FLOP/Byte，机器峰值 100 GFLOP/s、带宽 25 GB/s。AI* = 100/25 = 4，恰在拐点上，两者都可成为瓶颈；实测通常介于 min(100, 4×25) = 100 与实际带宽利用率之间。 例 3（Amdahl）：串行比例 S=5%，64 核最多加速 1/0.05 = 20×，实际 1/(0.05 + 0.95/64) = 13.9×。若负载不均导致某核多干 20% 的活（等效串行比例约 5%），加速比同样被限制在约 14×。 例 4（work-span）：一个算法的 T₁ = 1000 秒、T∞ = 10 秒 ⇒ 并行度 100；在 8 核上的时间下界 1000/8 + 10 = 135 秒 ⇒ 最高加速 7.4×，达不到 8×。

常考判断题：

• “我的代码在 8 核上加速 6.5×，所以 64 核上会加速 52×” ⇒ 错。要先看并行度（T₁/T∞）够不够，再看串行比例。
• “超线性加速不可能” ⇒ 错。当并行版本获得更好 cache 局部性（例如多个核的 cache 合起来能装下工作集）时可能出现 Speedup > P，此时应报告绝对时间与内存层次影响。
• “效率 100% 才算好” ⇒ 错。若 8 核用 4 倍功耗换 3 倍性能，从”性能/功耗”角度看是变差的；效率要与目标（时间还是能耗）一起谈。

A.7 高频考点五：硬件结构

流水线/ILP（L2）：

• 数据冒险（RAW/WAR/WAW）、控制冒险、结构冒险三类；旁路转发（forwarding）解决大部分 RAW；WAR/WAW 靠寄存器重命名消除。
• OoO 核心 = 前端（取指/译码/重命名）+ 指令窗口/ROB（乱序发射、顺序提交）+ 后端（发射端口/执行单元）。
• 提取性能的两把尺子：延迟界（依赖链长度 × 延迟）与吞吐界（操作数 / 发射宽度）；实际性能 = 两者中较慢者。
• 为什么 ILP 撞墙：窗口 O(W²) 的复杂度、单程序内可用并行度有限、频率受 C·V²·f 限制。

现代多核（L3）：

• 四个概念：多核（线程级并行）、SIMD（数据级并行）、内存延迟、内存带宽。
• 跑满机器的条件：核数 × 每核 IPC × 向量宽度 × SMT = 需要的独立工作份数（例如 16 核 × 8 宽 × 4 线程 = 512 份）。
• 延迟可以隐藏（预取、硬件多线程、ILP），带宽只能少访问。
• Roofline 判断：AI < P_peak/BW ⇒ 带宽受限，优化方向是”减少字节”（分块、融合、压缩数据），而不是”减少指令”。

互连网络（L10）：

• 拓扑比较维度：成本（连线数/端口数）、延迟（直径、平均跳数）、带宽（二分带宽）、可扩展性、路由复杂度。
• 关键数字感：crossbar 成本 O(N²)（N 端口）；ring 直径 N/2；mesh 直径 2(√N−1)；fat tree 二分带宽最优但成本高；Omega/蝶形网络用 log N 级换低成本但有阻塞。
• 流控三级粒度：message → packet → flit；store-and-forward 延迟 ∝ 跳数 × 包长，cut-through/wormhole 大幅降低；虚通道（VC）解决队头阻塞（HOL）。
• 公式：T = hops × (T_router + T_link)，理想情况下 T ≈ α + n/β。

缓存一致性（L11–13）：

• 为什么需要：私有 cache 复制数据 ⇒ 共享内存语义被破坏。
• 三条要求：写传播（write propagation）、写串行化（write serialization）。等价不变量：SWMR（Single-Writer-Multiple-Reader）+ Data-Value。
• 监听（snooping）：所有 cache 监听总线，state machine per cache line。
  • MSI：M（Modified，独占且脏）、S（Shared，多副本干净）、I（Invalid）。写 miss ⇒ BusRdX 使他人失效；读 miss ⇒ BusRd。
  • MESI：增加 E（Exclusive clean）：独占且与内存一致，可静默升级为 M 而不发总线事务 ⇒ 消除”先读后写私有数据”的 upgrade 开销。这是 MESI 相对 MSI 的核心收益。
  • MOESI / MESIF：O（Owned，脏但可共享，负责供给数据，免写回内存）；F（Forward，指定唯一响应者，减少重复响应）。
• 状态迁移必背表（MESI，本地请求）：

• 伪共享（false sharing）：不同变量落在同一条 cache line，被不同核写 ⇒ 一致性流量与真共享相同。这是”人为通信”，不是算法固有的。修法：padding 对齐、每线程本地副本 + 最后归并。
• 目录协议（L12）：每 cache line 一个目录项，含 P 个 presence bit + dirty bit；存储在 home node。读 miss 2 跳；写 miss 需要失效所有 sharer（2 + 2k 条消息、4 跳关键路径）。优化：limited pointer（只记录 k 个 sharer，溢出则回退广播）、sparse directory（只为在 cache 中的行保留项）、intervention forwarding（owner 直接回数据，降低关键路径）。
• 监听实现（L13）：原子总线 vs 拆分事务总线（请求/响应分离 + 请求表 + NACK 流控）；写回缓冲（write-back buffer）也要参与监听；取数死锁（fetch deadlock）、活锁、饥饿；写提交（commit）≠ 写完成（complete）；缓冲带来缓冲死锁，解法是请求队列与响应队列分离。

A.8 高频考点六：并行编程基础与性能优化（L7–L9）

四步法：Decomposition（分解）→ Assignment（分配）→ Orchestration（编排）→ Mapping（映射）。

工作分配：

• 静态 blocked：每个 worker 拿连续一块。优点：简单、零运行时开销、访存局部性好。缺点：负载不均时致命（Amdahl 放大）。
• 静态 interleaved：每个 worker 拿 i % P。负载更均匀、cache 利用可能更好，但访存不连续（共享内存下无妨，消息传递下会产生大量人为通信）。
• 动态：共享计数器 + chunk / 工作队列。开销 = 每次取号的原子操作与竞争（约 20–500 ns）。粒度公式：k ≥ c·P/t（避免抢任务的成本超过做任务）与 k ≤ M/(P)（保证足够多的任务数）。
• 工作窃取（work stealing）：每 worker 一个 deque，本地从尾部 push/pop，窃取者从头部拿；continuation stealing（先执行 spawn 的子任务，父任务留给窃取者）保证串行语义 + 减少队列操作。

通信代价：T(n) = T₀ + n/B，T₀ = overhead + occupancy + network delay。重叠加计算（overlap）可以把通信藏起来（如 MPI Isend/Irecv + Waitall）。 四类优化（讲义”四个 C”的实践）：

1. Blocking（分块）：提高重用距离内的时间局部性。
2. Fusion（循环融合）：把多遍遍历合成一遍，算术强度翻倍（例：三遍遍历 AI = 1/3 → 融合后 3/5）。
3. Sharing（共享）：把访问同一数据的任务放到同一执行单元（CUDA block 共享 L1/shared）。
4. 粒度（消息大小）：cache line（64 B）造成的边界浪费与伪共享。

竞争（contention）：热点资源吞吐上限 = 1/t_crit（t_crit = 该资源的关键路径延迟）。对策：复制资源（每线程私有 + 归并）、细粒度锁、分布式队列，或把串行化改写为”数据并行 + 排序”（讲义分桶例子）。

A.9 一页纸速记（可抄到 A4 上）

【三把尺子】
  Work T₁ / Span T∞ / 并行度 = T₁/T∞
  加速比 ≤ 1/(S + (1−S)/P)；上限 1/S
  AI = FLOP/Byte；P ≤ min(P_peak, AI×BW)；AI* = P_peak/BW

【三个墙】
  功耗墙 P ∝ C·V²·f  → 频率停涨 → 多核
  ILP 墙（窗口 O(W²)、程序并行度有限）→ 需要显式并行
  带宽墙（BW 增长 << 峰值算力）→ 大多数程序带宽受限

【硬件层次】
  寄存器 → L1(~4cyc) → L2(~12) → L3(~38) → DRAM(~200+)
  SIMD 宽度：SSE 4 / AVX2 8 / AVX-512 16 (float)
  GPU：SM = 4 sub-core × 32 lane = 128 lane；warp = 32 线程
  跑满机器所需独立工作 = 核数 × 每核IPC × 向量宽度 × SMT

【一致性】
  MSI：M(独占脏) S(共享净) I；MESI 加 E(独占净，可静默升级)
  读 miss → BusRd；写 miss → BusRdX；M 态被监听 → flush
  伪共享 = 不同变量同一条 line 被多核写 = 人为通信
  目录：presence bits + dirty bit；读 2 跳、写 2+2k 消息 4 跳
  limited pointer / sparse directory / intervention forwarding

【同步】
  锁性能看三段：acquire / waiting / release
  TAS → TTAS → ticket → 数组锁 / MCS 队列锁（O(1) 流量 + FIFO）
  屏障：集中式 sense-reversal（两次自旋变一次）/ 树形 lg P
  SC 太慢（写缓冲必须停顿）→ TSO/PSO/WO → 用 fence 与 acquire/release 补

A.10 自测 12 题

1. Q：为什么 MESI 中的 E 态能提升性能？ A：E 态表示”独占且干净”，处理器可以不发任何总线事务直接把 E 升级为 M（静默升级）。这消除了”先读后写一块私有数据”时的 BusRdX 事务；MSI 下同样的访问模式会产生读 miss（BusRd）+ 写 miss（BusRdX）两次事务。

2. Q：一个程序在 8 核上加速 7.2×，在 16 核上加速 7.4×。最可能的原因是什么？如何验证？ A：串行比例或并行度不足。由 Amdahl 得 S ≈ 1/7.4 − 小量 ≈ 13%，与 8 核的 7.2× 一致（1/(0.13+0.87/8) = 6.9）。验证方法：用 work-span 分析算出 T∞，或做强扩展/弱扩展曲线、最小二乘拟合 Karp-Flatt 度量 e = (1/S − 1/P)/(1 − 1/P)（e 接近常数说明并行开销固定占比）。

3. Q：为什么”给数组每个元素加 1”在 8 核上只有约 2× 加速？ A：算术强度极低（每 4 字节读 + 4 字节写只有 1 次加法 ⇒ 0.125 FLOP/Byte），远低于 ridge point，带宽受限；多核共享同一条内存总线，带宽不会随核数线性增长。实测受 STREAM 类上限约束。

4. Q：写出 ISPC 中交错分配与分块分配的差别，并说明哪个在 AVX2 上更快及原因。 A：见 A.3。交错更快，因为同一 gang 的并发访存地址连续，可合并为 packed load（vmovups 32 字节）；分块导致地址相距很远，需要 vgatherdps。

5. Q：__syncthreads() 能不能用来同步两个不同的 block？ A：不能。它只同步 block 内的线程。跨 block 同步必须用 kernel 分解（kernel 返回是全局隐式屏障）或原子操作 + 全局标志（但这有死锁风险，因为 block 不保证并发驻留）。

6. Q：目录协议的写 miss 需要多少条消息？关键路径多少跳？ A：2 + 2k 条消息（k = sharer 数）：请求→home、home 失效 k 个 sharer、k 个 ack 回 home、home 授权请求者（或用 intervention forwarding 让 owner 直接转发数据，关键路径降为 3 跳）。关键路径 4 跳（请求→home→sharer→home→请求者）。

7. Q：为什么”动态分配”不总是优于”静态分配”？ A：动态分配每次取任务都要做原子操作或加锁，有 20–500 ns 的竞争成本；若任务粒度太细（任务本身只有微秒级），调度开销会淹没收益。此外动态分配可能破坏局部性（任务与数据的亲和性丢失）。

8. Q：解释”延迟可以被隐藏，带宽不能”。 A：一次访存的延迟可以用其他独立工作填充（ILP、硬件多线程、预取），所以延迟是”能否找到并行工作”的问题；带宽是单位时间能搬运的字节数上限，是资源吞吐的物理上界，除了减少访问量（分块/融合/降低精度）无法绕过。

9. Q：什么是伪共享？为什么 padding 能解决它？ A：多个核写不同变量但落在同一条 64 字节 cache line 上，导致该 line 在所有写者之间反复迁移（ping-pong），一致性流量与真正共享同一变量时相同。padding 把每个线程的变量放进各自的 cache line（alignas(64)），使迁移消失。

10. Q：为什么”写提交（commit）”不等于”写完成（complete）”？ A：提交是指令从流水线退休（架构状态已更新），而完成是指该写在一致性协议中取得所有权并实际生效。x86 的 store buffer 使得本核看到的自己写的值早于其他核能看到的时间；这也是 TSO 的成因之一。

11. Q：FP32 矩阵乘在 V100（峰值 ~15.7 TFLOP/s、带宽 900 GB/s，ridge ≈ 17.4）上，若分块宽度 L=32，能达到多少性能？ A：分块后算术强度 AI = L/4 = 8 FLOP/Byte < 17.4 ⇒ 带宽受限，P ≤ 8 × 900 GB/s = 7.2 TFLOP/s（约为峰值的 46%）。若把 L 提到 64，AI = 16 ≈ 拐点；L = 128 时 AI = 32 > 17.4，转为算力受限。这也解释了为什么”分块尺寸”是 GPU GEMM 最重要的调优旋钮。

12. Q：8 核机器上某同步程序加速仅 3×，用 PMU 看到 BusRdX 事件随核数平方增长。最可能的两个原因？ A：① 伪共享（不同线程写同一 cache line 的不同字节）；② 集中式同步原语（全局自旋锁的 TAS 每次尝试都发 BusRdX，争用流量 O(P²)）。修法：alignas(64) 隔离 + 改 TTAS/ticket/MCS 锁或每线程私有累加。

第三部分：速查表与附录

速查表 A：并行编程模型速查表

选择决策树：

数据能否装进单节点内存？
├── 否 → MPI（+ 节点内 OpenMP/CUDA 混合）
└── 是 → 计算是否规则、数据量是否巨大（>10⁷ 独立元素）？
         ├── 是 → GPU（CUDA）；若只需 SIMD 加速 → ISPC
         └── 否 → 是循环级数据并行吗？
                  ├── 是 → OpenMP parallel for / reduction
                  └── 否 → 是递归分治或动态任务图吗？
                           ├── 是 → OpenMP task / Cilk / TBB（工作窃取）
                           └── 否 → 共享数据结构的细粒度操作？→ pthreads + 细粒度锁/无锁

跨模型对照：同一个”向量加”的六种写法

// 1. pthreads：手动分块 + join
for (t = 0; t < P; t++) pthread_create(&th[t], NULL, worker, &arg[t]);
for (t = 0; t < P; t++) pthread_join(th[t], NULL);

// 2. OpenMP：一个 pragma
#pragma omp parallel for schedule(static)
for (i = 0; i < N; i++) c[i] = a[i] + b[i];

// 3. ISPC：SPMD gang，编译器生成 SIMD
foreach (i = 0 ... N) c[i] = a[i] + b[i];

// 4. CUDA：成千上万轻量线程
__global__ void add(float *a, float *b, float *c, int n) {
    int i = blockIdx.x*blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}
add<<<(n+255)/256, 256>>>(a, b, c, n);

// 5. MPI：每个进程处理自己的切片，末尾可能需边界交换
int lo = rank * (n/P), hi = lo + n/P;
for (i = lo; i < hi; i++) c[i] = a[i] + b[i];
// 若 c 的邻居需要 a 的 halo，则 MPI_Sendrecv

// 6. C++17 并行 STL：最高层抽象，实现可能是线程池 + SIMD
std::transform(std::execution::par_unseq, a, a+N, b, c,
               [](float x, float y){ return x + y; });

速查表 B：性能公式速查表

B.1 可扩展性

B.2 内存与带宽

B.3 能耗与功率

B.4 计算峰值

峰值 FLOPS = 核数 × 时钟频率 × 每周期 FMA 数 × 2 × SIMD 宽度
  例：16 核 × 3.0e9 × 1 FMA × 2 × 8 (AVX2) = 768 GFLOP/s
  例：H100 SXM：132 SM × 128 FP32 lane × 2 × ~1.98 GHz ≈ 67 TFLOP/s (FP32)
                Tensor Core FP16（稀疏）可达约 2000 TFLOP/s

速查表 C：硬件架构参数速查表

C.1 存储层次典型延迟与容量

记忆要点：L1→L2→L3→DRAM 的数量级是 4 / 12 / 38 / 200+ 周期（约 3× 递增）。这些数字在考试里常用来做定量估算。

C.2 CPU vs GPU 结构对比

C.3 SIMD 宽度

C.4 一致性协议状态速查

C.5 内存一致性模型

速查表 D：优化技巧速查表

D.1 决策表：先量什么，再改什么

D.2 优化技巧清单（按影响量级排序）

内存/带宽类（通常收益最大）

1. 分块（blocking/tiling）：把工作集切到能装进 cache 的瓦片，循环顺序按瓦片嵌套。
2. 循环融合（fusion）：把多遍遍历合并成一遍，AI 直接翻倍。注意寄存器压力与可读性。
3. 重用（reuse）与数据布局：结构体数组 AoS → 数组结构体 SoA（利于向量化与合并访存）。
4. 消除人为通信：padding 对齐（alignas(64)）、每线程私有副本、减少 halo 交换频率。
5. 提高消息粒度：小消息合并成大消息（α 摊薄），例：把每次 8 B 的交换改成每次 4 KB。
6. 大页 / TLB 优化：madvise(MADV_HUGEPAGE)、mmap 2 MB 页；避免页步长（page-stride）访问大数组。

并行结构类

1. 提高并行度：把串行段并行化；用树形归约代替线性归约（span 从 O(P) 降到 O(log P)）。
2. 负载均衡：动态分配、工作窃取、长任务优先（LPT）、半静态自适应重划分。
3. 减少同步次数：合并屏障（diff[3] 轮转）、批量提交、每线程一次锁而非每次迭代。
4. overlap 通信与计算：MPI_Isend/Irecv + Waitall；CUDA stream 与 cudaMemcpyAsync。
5. 避免伪共享：写操作隔离到各自 cache line；读共享数据则无妨。

指令/SIMD 类

1. 向量化：-O3 -march=native、#pragma omp simd、ISPC foreach；检查生成汇编是否有 vfmadd*。
2. 消除分支发散：GPU 上让整个 warp 走同一路径；用算术或掩码代替分支（cmov）。
3. 合并访存（coalescing）：让相邻线程访问相邻地址（CUDA 尤其重要）；避免 stride 访问。
4. ILP 与展开：多条独立累加链打破延迟界；#pragma unroll。
5. 对齐：_mm_malloc(..., 32) / aligned_alloc(64, ...)，避免跨 cache line 的向量访问。

D.3 高水位实验（High Watermark）—— 判断”我的实现到底有多好”

方法：构造一个只能更快的”作弊版本”，测出物理上限，再看自己离它多远。

D.4 测量方法论 Checklist

• 用墙钟时间（omp_get_wtime、std::chrono::steady_clock），不用 clock()（那是 CPU 时间，多线程会累加）。
• warmup：先跑几次让 cache、TLB、线程池、频率进入稳态；CPU 频率缩放会把第一次运行变慢。
• 重复多次取中位数/最小值，报告分布而非单点。
• 防止死代码消除：把结果写进一个 volatile 全局或打印校验和。
• 基线要正确：串行版本必须是最优串行实现（不能用带锁的版本当基线）。
• 固定问题规模 vs 固定时间：明确报告是强扩展还是弱扩展。
• 关闭 turbo/绑核/控制变量：报告机器参数，避免跨机器混合比较。
• 用计数器而非感觉：perf stat -e cycles,instructions,cache-misses,LLC-load-misses；GPU 用 Nsight Compute 的 sm__throughput / dram__throughput 判断算力受限还是带宽受限。

速查表 E：同步原语与一致性协议速查表

E.1 锁的对比

E.2 屏障的对比

E.3 一致性协议状态迁移（MESI 核心表）

见 A.7 的状态迁移表。补充要点：

• 读未命中 S 态：发 BusRd，内存/其他 cache 供给数据，自己不写回。
• 写未命中：I → M 发 BusRdX；S → M 发 BusRdX（无效化其他副本）；E → M 静默。
• M 态被监听：BusRd ⇒ flush 数据并降到 S（或 F/O）；BusRdX ⇒ flush 并降到 I。
• 换出（eviction）：M 态换出必须写回内存（write-back）；O 态换出也需写回。

E.4 原子操作与硬件支持

速查表 F：常见性能陷阱速查表

附录 B：课程资源清单

B.1 官方链接

B.2 本地已下载材料（本工作目录）

cmu15418_data/
├── *.html                     课程网站全部公开页面（8 个）
├── lectures_data.json         结构化的逐讲资料记录（含公开性状态）
├── syllabus.pdf               课程大纲
├── doc_asst1_handout.pdf      Assignment 1 说明（13 页）
├── doc_CUDA-recitation.pdf    CUDA Recitation（39 页）
├── lectures/                  Fall 2026 公开讲义 PDF（21 份）+ 讲义逐页文本
├── extracted/                 38 份讲义/试卷的逐页文本抽取（含 slide 标记）
├── cs149_supp/                Stanford CS149 Fall 2025 公开讲义文本（18 份，补充材料）
├── supp/exams/                Spring 2022 练习题（4 份）+ 文本抽取
├── past/                      归档讲义下载尝试（实为 CMU 登录页，已记录为未公开）
└── notes/                     26 份逐讲学习笔记（本笔记主体）

B.3 相关课程与扩展阅读

结语：这门课的”最后一句话”

并行编程的本质不是”用更多核”，而是”用有限的硬件资源，以最小的数据移动，完成尽可能多的工作”。

• 硬件给了你：更多的执行单元（核 / SIMD lane / GPU lane）、多层次的存储、以及一套把并行性拼装起来的一致性机制。
• 软件要做的事：分解出足够多的独立工作、分配得足够均匀、编排得足够少的通信与同步、并映射到最合适的执行单元上。
• 判断标准永远只有三把尺子：Work/Span（并行度）、Amdahl（串行比例）、算术强度与 Roofline（瓶颈类型）。

祝复习顺利。