Lecture 9: Performance Optimization II: Locality, Communication, Contention

目录 · ← l8 · l10 →

Lecture 9: Performance Optimization II: Locality, Communication, Contention

1. 章节标题与概述

Lecture 9: Performance Optimization II: Locality, Communication, Contention（性能优化 II：局部性、通信与竞争）

• 本讲核心问题：上一讲（Performance Optimization I）解决的是怎样把工作合理地分配给 worker（静态/动态分配、粒度、工作窃取）；本讲解决另一半问题——怎样把通信的代价压到最低。讲义开篇即点明：”今天：最小化通信代价的策略。”要回答的具体问题是：一次通信到底贵在哪里（启动延迟、占用时间还是竞争）？哪些通信是算法固有的（inherent）、哪些是机器实现强加的（artifactual）？如何通过改变工作分配、改变数据布局、改变遍历顺序（blocking）、融合循环（fusion）、共享数据（sharing）、增大消息粒度来减少或隐藏通信？以及当多个 worker 同时争抢一个资源（锁、热点、带宽）时，性能为什么会崩塌？

• 涉及的主要硬件/软件机制：
  • 软件侧：消息传递的阻塞（synchronous/blocking）与非阻塞（asynchronous，Isend/Irecv + Wait/Test）语义、ghost cell（影子单元）数据复制、死锁回避的配对收发顺序、循环融合（loop fusion）、分块/瓦片遍历（blocking/tiling）、仿射工作分解（1D blocked / 1D interleaved / 2D blocked assignment）、细粒度锁（per-cell lock）、分布式工作队列（distributed work queues + work stealing）、数据并行的排序式分桶（counting sort + cell_starts/cell_ends）。
  • 硬件侧：存储层次（寄存器 → L1 → L2 → L3 → 本地内存 → 远端内存/网络）与 cache line 粒度（64 B）、缓存一致性带来的伪共享（false sharing）、cache line 在核间”乒乓”迁移、互连网络（Intel Sandy Bridge 起的 ring 互连：request/snoop/ack/data 四类环、4 个 L3 slice + system agent + graphics 共 6 个节点；SUN Niagara 2 的 crossbar；多路系统的 NUMA）、DRAM 带宽与延迟、memory-level parallelism（MLP）与硬件预取器。
  • 量化工具：延迟 vs 带宽的区分、流水线（pipelining）、通信代价模型 T(n) = T0 + n/B 与”overhead + occupancy + network delay”三段分解、算术强度（arithmetic intensity）与 communication-to-computation ratio、Roofline 模型、work-span（贪婪调度）模型、Amdahl 定律、Little’s law（带宽 = 并发度 / 延迟）、竞争与串行化排队。
• 在并行计算知识体系中的角色：本讲是第 7~8 讲”如何写出并行程序 / 如何分配工作”之后的性能收口：把”依赖 → 同步 → 通信”这条链从正确性推进到性能。它同时是后续若干讲的引子与上下文：
  • 第 11~13 讲（Snooping 一致性、目录协议、Snooping 多处理器设计）正是在回答”本讲提到的伪共享（两个处理器写同一条 cache line）在硬件上到底发生了什么”；
  • 第 10 讲（互连网络，Interconnection Networks）回答”网络拓扑如何决定 occupancy 与竞争”；
  • 第 15~16 讲（同步实现、细粒度同步与无锁编程；内存一致性在第 21 讲）回答”锁/原子操作的成本从何而来，为什么细粒度锁和树形结构能降低竞争”；
  • 第 14 讲（Performance Analysis / Profiling）与本讲”建立 high watermark、分清 compute / bandwidth / sync 瓶颈”的方法论直接衔接。 一句话概括本讲的地位：并行程序的可扩展性上限由通信与竞争决定，而通信与竞争既由算法（分解与分配）决定，也由机器（cache line、互连、锁实现）决定。
• 配套材料：
  • lectures/08_progperf2.pdf（对应抽取文本 extracted/08_progperf2.txt，共 63 页幻灯片）：已公开，位于 Fall 2026 公开讲义目录 https://www.cs.cmu.edu/~418/lectures/ 之下，可在公开网络直接下载。Fall 2026 日程表 https://www.cs.cmu.edu/~418/schedule.html 中本讲为 Lecture 9, Sep 14, Performance Optimization II，日程表已给出 slides 链接（视频链接被注释隐藏）。
    • 讲义首页写作 “Lecture 7: Performance Optimization Part II: Locality, Communication, and Contention”，页脚学期字样为 Fall 2025：这是讲义沿用历史版本的正常现象，不代表日程表出错（Fall 2026 的讲次编号顺延为第 9 讲）。
    • 上一讲 lectures/07_progperf1.pdf（Performance Optimization I，Fall 2026 为 Lecture 8, Sep 11）是本讲的直接前置：那里讲工作分配（静态/动态、粒度、Cilk 式 continuation stealing 工作窃取），本讲讲通信、局部性与竞争。
  • cs149_supp/perfopt2.txt：已公开的姊妹课程材料 —— Stanford CS149 Fall 2025 Lecture 6 “Performance Optimization Part II: Locality, Communication, and Contention”（68 页）。它与 CMU 版本共享同一套骨架与案例（Culler/Singh/Gupta 的网格求解器与 ghost cell、α-β 通信模型、inherent vs artifactual communication、office hours 竞争类比、分布式工作队列、Roofline、high watermark），但多出四块内容：共享地址空间硬件的实现（ring/crossbar/NUMA）、性能分析方法论（Roofline、perf counter、PCM）、固定问题规模测速的陷阱（SGI Origin 2000 上的 258×258 网格）与超线性加速。本笔记在使用这些额外内容时会显式标注来源为 CS149 版本。
  • 其它讲座的公开状态（与本讲无关但同属 Fall 2026）：Performance Analysis / Profiling（日程表第 14 讲）、Transactional Memory（第 17 讲）、AI in System Design（第 24 讲）在 Fall 2026 尚未给出 slides 链接，属未发布；历史学期对应 PDF 位于 /afs/cs/academic/class/15418-*/public/ 之下，需要 CMU 登录，属未公开。
  • 录像（Panopto/YouTube）在 Fall 2026 日程表中被注释隐藏，属 未发布；Ed 讨论区、Autolab、Canvas 均需登录，非公开。
  • Fall 2026 授课教师为 Brian Railing 与 Dimitrios Skarlatos；课程由 Kayvon Fatahalian 创建。

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

2.1 延迟 vs 吞吐量：一切通信优化的出发点

• 定义与目的：
  • 延迟（latency）：完成一次操作所需的时间。讲义给的定义是”The amount of time needed for an operation to complete”，例子：一次 cache miss 的 load 有 ~200 个周期的延迟；一个 packet 从本机到 Google 要 20 ms；在 Piazza 上提问 10 分钟得到回复。
  • 带宽（bandwidth / throughput）：操作被完成的速率。例子：内存可以 25 GB/s 的速率向处理器供数；一条链路每秒能发 1000 万条消息；助教每天回答 50 个问题。
  • 区分二者是通信优化的第一步：延迟决定”一次通信要等多久”，带宽决定”单位时间能传多少”；二者需要完全不同的优化手段（降延迟 vs 提吞吐）。
• 直观解释（”它是什么？”）：讲义用”从旧金山开车去匹兹堡”作类比。距离约 4000 km，车速 100 km/h ⇒ 一个人的延迟 = 40 小时。如果高速公路上车与车之间保持 1 km 间距，那么吞吐量 = 100 人/小时（每 1/100 小时过一辆车）。
  • 提高吞吐量有两条路：① 开快点（150 km/h ⇒ 150 人/小时）；② 多修车道（3 车道 ⇒ 300 人/小时）。前者相当于提升”单个流的速度”（提高时钟频率、降低单次延迟），后者相当于复制资源（更多执行单元、更多链路）。
  • 极端情形：如果高速公路上同一时刻只允许一辆车，那么吞吐量退化为 1 / 延迟 = 1/40 人/小时。这就是”没有重叠的通信”——最糟糕的情形，也是本讲要消灭的情形。
  • 另一个更贴近生活的类比：同时把”车距”减半（1 km → 500 m）可以把吞吐量翻倍到 200 人/小时，效果等价于在 1 km 间距下把车速提到 200 km/h。这就是流水线（pipelining）的核心思想：把一个大任务切成更小的子任务，让不同的”车”（不同的阶段）并行处理不同部分，而不是让一辆车独自跑完全程。
• 架构/机制图解（流水线的时空图）：

  (a) 非流水线 / 串行执行：吞吐量 = 1 / 延迟
      时刻:   0h        1h        2h        3h        4h        5h        6h
      load#1  [====== 洗涤 45min ======][==== 烘干 60min ====][折 15m]  <- 2 小时/批
      load#2                                                      [====== ...       <- 必须等前一批做完
      吞吐量 = 1 批 / 2 小时 = 0.5 批/小时           延迟(1 批) = 2 小时

  (b) 复制资源（两台洗衣机 + 两台烘干机 + 一个朋友）：延迟不变，吞吐量 x2
      load#1  [== 洗 ==][== 烘 ==][折]
      load#2  [== 洗 ==][== 烘 ==][折]        吞吐量 = 2 批 / 2 小时 = 1 批/小时
                                                资源数量也 x2（两套设备 + 两个人）

  (c) 流水线：一套设备，把"批次"错开，资源始终不空闲
      时刻:   0h      1h      2h      3h      4h      5h
      洗衣机  [L1洗][L2洗][L3洗][L4洗][L5洗] ...
      烘干机        [L1烘][L2烘][L3烘][L4烘] ...      <- 烘干 60min 是"最慢环节"
      学生          [L1折][L2折][L3折][L4折] ...      <- occupancy = 1 小时
      吞吐量 = 1 批 / 小时（= 1 / 最慢环节占用的时间）
      延迟   = 2 小时（单批依然要 洗45 + 烘60 + 折15）

  (d) 四段指令流水线 IF -> D -> EX -> WB：延迟 = 4 周期，吞吐量 = 1 指令/周期
      clk:      1     2     3     4     5     6     7
      instr0   IF    D    EX    WB
      instr1         IF    D    EX    WB
      instr2               IF    D    EX    WB
      instr3                     IF    D    EX    WB
      (真实机器更长：讲义提到 Intel Core i7 的流水线是变长的 ~15-20 级)
      前提：相邻指令之间存在依赖时，必须由旁路/停顿/乱序来保证正确性。

• 性能特征与关键结论：
  • 流水线的延迟不变（单批仍是 2 小时；单条指令仍是 4 周期），但吞吐量被”最慢一级的占用时间”（occupancy）锁定：稳态下 吞吐量 = 1 / occupancy。
  • 流水线不需要复制执行资源（只有一台洗衣机、一台烘干机），这是它比”复制资源”更便宜的地方；代价是单次操作的延迟不再下降，而且需要额外的并发（讲义后文：”requires additional concurrency in application — more concurrency than number of execution units”）。
  • 对通信而言：“修更多车道”= 提高链路带宽/增加链路条数；”开快点”= 降低单次传输延迟；”流水线”= 让多笔通信在时间上重叠。本讲后面的全部技术都可以归到这三类里。

2.2 通信代价模型：T(n) = T0 + n/B

• 定义与目的：给出一个能预测通信成本的模型，从而知道优化应该往哪里使力。讲义先给最简单的”非流水线通信”模型：
  • T(n)：一次传输的总延迟；T0：启动延迟（start-up latency），即第一个 bit 到达目的地所需的时间；n：传输的字节数；B：链路传输速率（带宽）。
  • T(n) = T0 + n / B；有效带宽（effective bandwidth） = n / T(n)。有效带宽依赖于传输大小：小消息被启动延迟摊薄得很惨，大消息才能逼近 B。
  • 更一般的模型把一次通信拆成三部分：总通信时间 = overhead（开销）+ occupancy（占用）+ network delay（网络延迟）。其中 overhead 是处理器花在通信上的时间（发起调用、拷贝到网络缓冲）；occupancy 是数据穿过系统最慢部件所需的时间；network delay 是其余所有时间。
  • cost（代价）= 通信时间 − overlap（重叠部分）：重叠掉的部分不产生代价。这正是非阻塞通信与流水线的意义。
• 直观解释（”它是什么？”）：类比寄快递。overhead = 你在家打包、填单、交给快递员的时间（你自己被占住了）；occupancy = 包裹在传送带上/货车上的时间（系统资源被占住了，宽窄由最慢那段路决定）；network delay = 中转分拣、堵车。如果你有 10 个包裹要寄，一次打包寄一个（非流水线）与一次打包 10 个一起寄（大消息/合并消息）的总时间差别巨大——这就是 T0 的摊薄效应。而你一边等快递一边继续干活，就是 overlap。
• 架构/机制图解（三段分解与重叠）：

  发送方                                                          接收方
  ┌────────┐   ┌─────────────┐   ┌──────────────┐   ┌──────────┐
  │ SEND() │──►│ 网络缓冲/链路│──►│ 慢链路(小 B) │──►│  拷贝到  │
  │ 调用+  │   │   传输       │   │ 快链路(大 B) │   │ 接收缓冲 │
  │ memcpy │   └─────────────┘   └──────────────┘   └──────────┘
  └────────┘
   <-- overhead -->|<--net delay-->|<---- occupancy: T0 + n/B(慢) ---->|<-overhead->

  时间线（多次发送，非流水线 vs 流水线）：
   非流水线（每笔都等上一笔完成，吞吐 = 1/T(n)）:
     msg0 |==T0==|====n/B====|
     msg1                     |==T0==|====n/B====|
     msg2                                          |==T0==|====n/B====|
   流水线（本笔的 n/B 与下一笔的 T0 重叠，吞吐 -> 1/occupancy）:
     msg0 |==T0==|====n/B====|
     msg1        |==T0==|====n/B====|
     msg2               |==T0==|====n/B====|
              ↑ 稳态下 B 一直满负荷：吞吐 = 1/occupancy

   注意：若网络缓冲最多只能放 2 条消息（讲义假设），突发的发送速率
   快于 1/occupancy 时，发送方会被"缓冲已满"阻塞 —— 流水线深度受缓冲限制。

• 性能特征与关键结论：
  • 有效带宽随消息变大而提高（n/T(n) = n/(T0+n/B) → B）：这直接推导出讲义的第一条优化准则——“send fewer messages, make messages larger”（减少消息条数、增大消息体积、把许多小消息合并成一个大消息）。
  • 定量地：要让有效带宽达到 B 的 50%，需要 n/B ≥ T0，即 n ≥ T0·B。以 T0 = 2 μs、B = 6.4 GB/s 为例，n ≥ 12.8 KB——小于 12.8 KB 的消息，一半以上的时间都花在启动上。
  • occupancy 决定稳态通信速率（msg/sec = 1/occupancy），而 occupancy 由最慢部件决定：一个”快链路 + 慢链路”的组合，最终速率由慢链路决定。
  • 只有当应用有额外并发（比执行单元更多的并发）时，overlap 才可能；否则”一边通信一边计算”只是口号。

2.3 消息传递与 ghost cell：把通信显式化

• 定义与目的：消息传递模型（message passing）中每个线程/进程拥有自己的私有地址空间，线程之间只能通过 send/recv 交换数据；send 指定接收者、被发送的缓冲区与可选的消息标识（tag），recv 指定发送者、接收缓冲区与 tag。它的价值在于：通信被显式写进代码，程序员被迫面对”谁拥有数据、什么时候需要谁的数据”，因此程序结构天然更接近可扩展的并行程序（集群/超算的编程模型，可用 Infiniband 之类的网络把商品机器连成大型并行机）。讲义之所以在性能优化课上先复习它，正是因为它是讨论通信代价最干净的载体。
• 直观解释（”它是什么？”）：类比每家自备仓库 + 寄信。共享地址空间是”大家共用一个白板”，写上去别人立刻能看见；消息传递是”每人有自己的仓库，想知道别人的数据必须写信索取”。于是出现一个关键概念：ghost cell（影子单元）——为了算自己负责的边缘单元，必须在本地复制一份邻居的数据，并明确”这份数据的所有权属于别的线程”。这就像你要用邻居家的温度计读数，只能抄一份贴在自家墙上，而抄来的那份可能过期。
• 架构/机制图解（SPMD 网格求解器 + ghost 行交换）：

   1D 行分块 (1D blocked assignment)：N x N 网格切成 P 份，每份 rows_per_thread = N/P 行
   每个线程的本地数组 = (rows_per_thread + 2) 行 x (N + 2) 列   ← 上下各 1 行 ghost，左右各 1 列 halo

        线程1 地址空间            线程2 地址空间            线程3 地址空间
     ┌──────────────────┐    ┌──────────────────┐    ┌──────────────────┐
     │      ...         │    │  ghost 上行 (白) │◄───┼─ send row ───┐   │
     ├──────────────────┤    ├──────────────────┤    ├──────────────┼───┤
     │  内部第 1 行(红) │───►│  内部第 1 行     │    │              │   │
     │  内部第 2 行     │    │  内部第 2 行     │    │              │   │
     │      ...         │    │      ...         │    │              │   │
     │  内部第 R 行(红) │───►│  内部第 R 行     │    │              │   │
     ├──────────────────┤    ├──────────────────┤    ├──────────────┼───┤
     │                  │    │  ghost 下行 (白) │◄───┼─ send row ───┘   │
     └──────────────────┘    └──────────────────┘    └──────────────────┘
        ▲ 竖直方向相邻的线程互发"边界行"；水平方向靠本地 halo 列解决

   Solve() 每轮循环的结构（slide 给出的伪代码）：
     while (!done) {
       ① 收发 ghost 行：  if (tid!=0)      send(&localA[1,0],  ..., tid-1, MSG_ROW);
                          if (tid!=P-1)    send(&localA[R,0],  ..., tid+1, MSG_ROW);
                          if (tid!=0)      recv(&localA[0,0],  ..., tid-1, MSG_ROW);
                          if (tid!=P-1)    recv(&localA[R+1,0],..., tid+1, MSG_ROW);
       ② 计算：5 点平均  localA[i,j] = 0.2*(上+下+左+右+自己)，同时累加 |Δ|
       ③ 归约：所有线程把 my_diff 发给线程 0；线程 0 判断 my_diff/(N*N) < TOLERANCE
               并把 done 广播回所有线程（用消息实现"屏障/标志位"）
     }

  ⚠ 阻塞式 send/recv 的死锁：所有线程"先 send 再 recv"会互相等待
      T0: send->T1 ...阻塞...           T1: send->T0 ...阻塞...
      T0: ...等着收 T1 ...              T1: ...等着收 T0 ...
      => 死锁（谁都不往下走）
      修法（slide 给出的版本）：奇偶分角色、收发顺序配对
      if (tid % 2 == 0) { sendDown(); recvDown(); sendUp(); recvUp(); }
      else              { recvUp();   sendUp();   recvDown(); sendDown(); }
      T0: [send][recv][send][recv]      T1: [recv][send][recv][send]
          ↑ 双向的收发在时间上交错，形成"配对"从而不会彼此死等
      另一条路：使用非阻塞 Isend/Irecv —— 把所有收发一次性投递出去再 Waitall

• 阻塞 vs 非阻塞（讲义的两张对照图）：

• 性能特征与关键结论：
  • 消息传递里通信与同步是同一件事：互斥、屏障、标志位都可以用消息实现（harness 里”用消息实现屏障”是常见练习）。
  • 批量传输（bulk transfer）是天然的优化：一次发一整行（N+2 个 float）比逐元素发 N+2 条消息快几个数量级——这是 T0 摊薄效应的直接应用。
  • 数组索引是相对本地地址空间的，任何”跨进程”的访问都必须显式翻译为消息；这消除了共享内存编程里”看不见的通信”。

2.4 把并行系统看作”扩展的存储层次”

• 定义与目的：讲义要求把”通信”这个概念一般化：不只是机器之间的消息，还包括处理器与 cache 之间、与本地内存之间、与远端内存之间的所有数据搬运。整个并行系统因此可以画成一张扩展的存储层次（extended memory hierarchy）：越靠近处理器，延迟越低、带宽越高、容量越小；任何一次”本地没法满足”的访问都会触发与下一级之间的通信。
• 直观解释（”它是什么？”）：类比一个人的办公桌：常用的笔和便签在手边（寄存器/L1）；当天要用的资料在抽屉（L2/L3）；不常用但要留档的在楼下档案室（本地内存）；别人部门的档案要打电话去要（远端内存，1 个网络跳）；上级单位的档案要层层转达（N 个网络跳）。你每多伸手一次就多一份往返时间，所以”把相关材料先搬到手边再干活”（局部性）永远是最有效的优化。
• 架构/机制图解（扩展存储层次 + 典型数量级）：

                        ┌───────────────┐
                        │   寄存器 Reg   │   ~1 cycle,  几十个,  数百 B/cycle
                        └───────┬───────┘
                                │
                        ┌───────▼───────┐
                        │  本地 L1 (32K)│   ~4 cycle,   64 B/line
                        └───────┬───────┘
                                │
                        ┌───────▼───────┐
                        │  本地 L2 (256K│   ~12-20 cycle
                        │  或 512K)     │
                        └───────┬───────┘
                                │
                ┌───────────────▼───────────────┐
                │   L3 / 共享 cache (数十 MB)   │   ~40-80 cycle
                └───────────────┬───────────────┘
                                │
                ┌───────────────▼───────────────┐
                │       本地内存 DRAM           │   ~200 cycle(60-100 ns), 带宽 B_D
                └───────────────┬───────────────┘
                                │
                ┌───────────────▼───────────────┐
                │  别的核的 L2 / L3 slice       │   与远端内存同量级
                └───────────────┬───────────────┘
                                │
                ┌───────────────▼───────────────┐
                │  远端内存（1 个网络跳）        │   100 ns ~ 1 μs, 带宽低得多
                └───────────────┬───────────────┘
                                │
                ┌───────────────▼───────────────┐
                │  远端内存（N 个网络跳）        │   延迟随跳数线性增长
                └───────────────────────────────┘
      从下往上：延迟更低、带宽更高、容量更小；
      从上往下：延迟更高、带宽更低、容量更大。
      一次"未命中"就是一次通信 —— 所以"管理局部性"在**每一层**都重要。

• 性能特征与关键结论：
  • 层次结构决定了”通信”无处不在：一个 load 指令可能引发 L1 未命中 → L2 → L3 → 本地内存 → 甚至远端内存的连锁通信。
  • 由此得到本讲的中心命题：优化局部性 = 优化通信 = 优化性能，而且这个判断在任何层次上都成立（register↔L1、L1↔L2、core↔DRAM、node↔node）。
  • CS149 版本给出的硬件实例（说明”共享地址空间抽象的实现有多复杂”）：Intel Sandy Bridge 起的 ring 互连（四条环分别传 request/snoop/ack/data，data 环 32 B；6 个节点 = 4 个 2 MB 的 L3 slice + system agent + graphics；每个 L3 bank 接两次环；3.4 GHz 下”每核访问自己的本地 L3 slice”时理论峰值带宽约 435 GB/s）；SUN Niagara 2 的 crossbar（全连接，面积约等于一个核心）；多路系统的 NUMA（不同核访问同一地址的延迟与带宽都可能不同，即使在单路系统里，不同 L3 slice 与核的距离也不同）。

2.5 固有通信 vs 人为通信（inherent vs artifactual communication）

• 定义与目的：这是本讲最重要的分类。
  • 固有通信（inherent communication）：给定工作分配方式，算法必须在处理器之间搬运的信息量（假设 cache 容量无限、传输粒度无限小）。例：消息传递求解器里的 ghost 行就是固有通信。它由”分解 + 分配”决定，因此程序员可以通过改进分配来降低它。
  • 人为通信（artifactual communication）：其余所有通信，源自系统的实现细节（无限容量 cache 假设、最小传输粒度假设不成立）。它往往和固有通信一样重要。
  • 人为通信的四个典型来源（讲义原文）：
    1. 系统有最小传输粒度：程序只要 1 个 4 字节 float，硬件却必须搬一整条 64 字节 cache line ⇒ 16 倍的多余通信。
    2. 系统操作规则导致的”多余”通信：程序连续写 16 个 4 字节 float（正好一条 64 B line），硬件却先把这条 line 从内存读进来（write-allocate），然后整条被覆盖，最后再写回 ⇒ 2 倍开销（读其实完全没必要）。
    3. 数据在分布式内存中放置不当：数据不在”访问它最多的处理器”附近。
    4. 复制容量有限：cache 太小留不住数据，同一份数据被反复搬运（capacity miss）。
• 直观解释（”它是什么？”）：类比搬家。固有通信 = 你确实必须搬走的家具（不管你用几辆车、多大的箱子，这些家具一定要运）。人为通信 = 因为你只买到了固定尺寸的纸箱（最小粒度），所以你必须为一只杯子用掉一整个大箱子（4 B → 64 B）；因为搬运规则要求先把箱子搬进屋再决定要不要（write-allocate），所以空箱子也要搬一趟（2 倍）；因为仓库太小放不下你已经搬来的东西，所以同一件家具你搬了两遍（capacity miss）。
• 架构/机制图解（工作分配如何决定固有通信量）：

   N x N 网格，P 个处理器，5 点模板（每轮读上下左右 4 个邻居）

   ① 1D blocked（连续行分块，P=4）          ② 1D interleaved（隔 P 行取一行，P=4）
      ┌────┬────┬────┬────┐                    ┌────┬────┬────┬────┐
      │ P0 │ P1 │ P2 │ P3 │  ← 每个处理器      │ P0 │ P1 │ P2 │ P3 │  ← 每个处理器
      ├────┼────┼────┼────┤     拿到 N/P 行    ├────┼────┼────┼────┤     拿到 N/P 行
      │ P0 │ P1 │ P2 │ P3 │     连续的网格      │ P0 │ P1 │ P2 │ P3 │     但它们被"打散"
      ├────┼────┼────┼────┤                    ├────┼────┼────┼────┤     到整个网格上
      │ P0 │ P1 │ P2 │ P3 │                    │ P0 │ P1 │ P2 │ P3 │
      ├────┼────┼────┼────┤                    ├────┼────┼────┼────┤
      │ P0 │ P1 │ P2 │ P3 │                    │ P0 │ P1 │ P2 │ P3 │
      └────┴────┴────┴────┘                    └────┴────┴────┴────┘
      计算/处理器 = N²/P                        计算/处理器 = N²/P
      通信/处理器 ≈ 2N  （上下两条完整边界行）   通信/处理器 ≈ 2N·(N/P) = 2N²/P
        （自己的行彼此相邻，接收到的行可复用）      （自己的行彼此间隔 P 行，
                                                    每条自己的行都要向邻居另要 2 行）
      算通比 = N/(2P)  ← 随 N 增大而变好          算通比 = 1/2  ← 与 N、P 无关，永远差

   ③ 2D blocked（√P x √P 的块，P=16）
      ┌────┬────┬────┬────┐      计算/处理器 = N²/P
      │ P0 │ P1 │ P2 │ P3 │      通信/处理器 ≈ 4·(N/√P)     ← 四条边，每条 √P 个元素... 
      ├────┼────┼────┼────┤                                  具体为每条边长 N/√P
      │ P4 │ P5 │ P6 │ P7 │      算通比 = N/(4√P)  ⇒ 算术强度随 N 增大而提高，
      ├────┼────┼────┼────┤                       且随 P 增大只是"次线性"变差
      │ P8 │ P9 │P10 │P11 │      （讲义的结论：communication costs increase
      ├────┼────┼────┼────┤        sub-linearly with P；分配方式捕获了算法的
      │P12 │P13 │P14 │P15 │        二维局部性）
      └────┴────┴────┴────┘

• 定量对照表（N×N 网格、P 个处理器、单位：元素）：

约定：本表把算术强度（arithmetic intensity）定义为”每搬运一个元素所对应的计算元素数”（越大越好）。讲义在 2D blocked 一页里同时写了 N/(4√P) 与 4√P/N，后者是它的倒数（即 communication-to-computation ratio，越小越好），二者只是定义方向不同。

• 性能特征与关键结论：
  • 算术强度 = 1 / communication-to-computation ratio。讲义原话：”I find arithmetic intensity a more intuitive quantity, since higher is better.” 现代并行处理器的计算能力与可用带宽之比很高，因此必须有高的算术强度才能高效利用机器（回想第 3 讲”逐元素向量乘法”的例子：几乎全部时间在等数据）。
  • “计算机比内存快得越来越多”是算术强度重要的根本原因：如果算力每两年翻倍而带宽只翻 1.5 倍，那么任何算术强度不变的程序都会自动变得越来越受带宽限制。
  • 好的分配决策可以降低固有通信，这是程序员在白板上就能做的优化（不需要改硬件）；而人为通信要靠数据布局与代码重排来消灭（见 2.7）。

2.6 四个 C：从 working set 视角看通信

• 定义与目的：讲义把经典的”三个 C”扩展为”四个 C”，把并行系统里的通信也纳入 cache miss 的分类：
  1. Cold miss（冷缺失）：第一次访问该数据。在串行程序里不可避免（”compulsory miss”）。对并行程序，它对应”第一次拿到工作集”（first working set）。
  2. Capacity miss（容量缺失）：工作集比 cache 大。可通过增大 cache 或改变分块/遍历顺序来减少。
  3. Conflict miss（冲突缺失）：由 cache 的管理策略（组相联度）造成。可通过提高相联度或改变数据访问模式（padding、分块布局）减少。
  4. Communication miss（通信缺失，本讲新增）：由并行系统中的固有通信或人为通信造成。
• 直观解释（”它是什么？”）：把 cache 想成书桌，内存是书架。Cold miss = 这本书你第一次去书架上取（不可避免）。Capacity miss = 你要同时摊开 20 本书，但桌子只放得下 5 本，于是来回跑。Conflict miss = 桌上明明有空位，但书架的编号规则让你的 6 本书只能落在同一个格子，只能互相挤出去。Communication miss = 这本书在隔壁同学的桌上，你必须跟他要——而且如果你每次都要，说明你的”工作分配”没把相关的书放在同一张桌上。
• 架构/机制图解（working set 视角下的数据流量）：

   数据流量
     ▲
     │  ┌──── 冷缺失 (Cold misses)：第一次触碰工作集，串行程序也躲不掉
     │  │
     │  ├──── 固有通信 (Inherent communication)：并行算法/分配方式决定的下界
     │  │
     │  └──── cache 容量造成的流量 (capacity + conflict)：
     │          工作集超出容量 → 反复搬运（人为通信的主体）
     │
     └──────────────────────────────────────────────► 层次容量的增加
        "first working set"      "second working set"
        (能装下第一个工作集)      (连更大的工作集也装得下)

   架构师的问题：给定这个负载，这一级该造多大容量？   ← 答案就在这张曲线的"拐点"上
   程序员的问题：我的程序落在这条曲线的哪一段？我能把它推到更低的那一段吗？
   注意：这张图在并行系统的**任何一层**都成立（register/L1/L2/L3/本地内存/远端内存/网络）

• 性能特征与关键结论：
  • 曲线有三段平台：容量太小 → 数据流量随容量增加陡降（capacity/conflict 主导）；容量足够后趋于平缓（只剩 cold miss + 固有通信）。
  • CS149 版本补充：”这张图看起来很熟悉吧？”——它就是经典的 memory mountain / working set 曲线（回想 15-213 cache lab）。
  • 工程含义：在”曲线仍然陡峭”的区间加容量收益最大；而如果程序已经落在平缓段，再大的 cache 也救不了它——这时应该改算法/分配方式（降低固有通信）或改布局（降低人为通信）。

2.7 降低通信的四种技术：blocking、fusion、sharing、granularity

2.7.1 时间局部性：改变遍历顺序（blocking / tiling）

• 定义与目的：分块（blocking）是”重排计算顺序以减少 capacity miss”的通用技术：把一个大的二维（或多维）迭代空间切成能整块装进 cache 的瓦片，在瓦片内部把所有能用到的数据一次用完，再去下一块。
• 直观解释（”它是什么？”）：类比做菜。朴素做法：为了做 3 道菜，先跑一趟超市买 A 的料，做 A；再跑一趟买 B 的料，做 B；再跑一趟买 C 的料（同一批调料被反复买回家）。分块做法：先想清楚要做的 3 道菜共用什么调料，一次买齐、在台面上按顺序做完（调料只买一次，且在桌上留住）。
• 架构/机制图解（讲义给出的具体模型：cache line = 4 个网格元素，cache 容量 = 24 个元素 = 6 条 line）：

   行主序遍历（row-major traversal）：
     N x N 网格，处理"红色元素"所在的第 i 行时，要读 i-1、i、i+1 三行。
     由于一行太长（远超 24 个元素），处理到下一行时，上一行的数据已被挤出 cache。

     cache 里最终留下的（蓝）= 第 i 行与 i+1 行的部分 + 第 i+2 行的开头
     ┌──────────────────────────────────────────────┐  ← 第 i 行 (已处理)
     │██████████████████████████                    │     蓝色 = 还在 cache 里
     ├──────────────────────────────────────────────┤  ← 第 i+1 行 (已处理)
     │██████████████                                │
     ├──────────────────────────────────────────────┤  ← 第 i+2 行 (即将处理)
     │██                                            │
     └──────────────────────────────────────────────┘
       ⇒ 每写 4 个输出元素，就要从内存加载 3 条 line（3 lines / 4 elements）
         其中 (0,2)、(1,1) 等元素"之前明明访问过"，但处理第 2 行时已经不在 cache 里
         —— 这是 **capacity miss**（工作集 = 3 行 > cache 容量）

   分块遍历（blocked iteration order）：
     把网格切成 8 列 x 若干行的瓦片（第一行块示意）：
     ┌────────┬────────┬────────┬────────┐
     │  瓦片0 │  瓦片1 │  瓦片2 │  瓦片3 │   ← 依次处理列方向上的瓦片
     │ [8 列] │        │        │        │      每个瓦片内的 3 行 x 8 列 = 24 个元素
     ├────────┼────────┼────────┼────────┤      正好塞进 cache（6 条 line）
     │  瓦片4 │  瓦片5 │  瓦片6 │  瓦片7 │
     ├────────┼────────┼────────┼────────┤
     │   ...  │        │        │        │
     └────────┴────────┴────────┴────────┘
       ⇒ 现在每加载 3 条 line 可以覆盖 8 个输出元素（3 lines / 8 elements）
       （CS149 版本在它自己的 cache 模型下给出"2 条 line / 6 个元素"；
         两者都是同一结论：**分块把每条 line 的复用次数从 ~1.3 次提到 ~2.7 次**）

• 性能特征与关键结论：
  • 判据是重用距离（reuse distance）：如果”两次使用同一数据之间要访问的数据量”超过了 cache 容量，就必然发生 capacity miss。分块的本质是把重用距离压到 cache 容量以下。
  • 分块的代价：循环开销增加、代码更复杂、边界处理更啰嗦；且块大小需要针对 cache 容量调参（太大装不下、太小则循环开销占比高）。
  • 讲义用这张图把 15-213 的 cache lab 与并行网格求解器联系起来：同一套局部性原理，串行并行通吃。

2.7.2 时间局部性：循环融合（loop fusion）

• 定义与目的：循环融合把多个”遍历同一批数组”的循环合并成一个循环，从而不再需要中间临时数组，把多次”读—写—再读”的内存往返压缩成”一次读完所有输入、一次写出结果”。
• 直观解释（”它是什么？”）：类比洗菜—切菜—炒菜。不融合 = 把所有菜洗一遍放进盆 A，再把盆 A 的菜全部切一遍放进盆 B，再全部下锅（每道工序都要把整批菜端来端去，还要多两个盆）。融合 = 拿起一棵菜，洗完立刻切、切完立刻下锅（每个菜只经过一次手，不需要盆）。
• 机制图解与算术强度（讲义原例）：

   不融合（模块化写法，类似 numpy 的数组运算）：
     tmp1[i] = A[i] + B[i];     // 读 A,B 写 tmp1      (每次数学运算: 2 load + 1 store)
     tmp2[i] = tmp1[i] * C[i];  // 读 tmp1,C 写 tmp2   (每次数学运算: 2 load + 1 store)
     E[i]    = tmp2[i] + D[i];  // 读 tmp2,D 写 E      (每次数学运算: 2 load + 1 store)
     总流量(每元素) = 3 load + 3 store = 9 次 4B 搬运
     算术强度 = 3 个数学运算 / 9 次搬运 = **1/3**
     —— 而且 tmp1、tmp2 是两个"只用一次"的中间数组（可能各占 64 MB！）

   融合（一次遍历）：
     E[i] = D[i] + (A[i] + B[i]) * C[i];   // 读 A,B,C,D 写 E
     总流量(每元素) = 4 load + 1 store = 5 次 4B 搬运
     算术强度 = 3 个数学运算 / 5 次搬运 = **3/5**
     ⇒ 内存流量降到 5/9 ≈ 56%（理论加速比上限 1/0.56 ≈ 1.8 倍）

   内存流量的直观对比：
     不融合:  A ──┐         ┌── C        D            (2 个临时数组来回过手)
                  ├─[tmp1]─┤
             B ──┘         └─[tmp2]── E
             流量: 读A+B 写tmp1 | 读tmp1+C 写tmp2 | 读tmp2+D 写E   = 9 趟
     融合:    A,B,C,D ──[一次循环体]── E
             流量: 读 A,B,C,D 写 E                                  = 5 趟

• 性能特征与关键结论：
  • 融合的收益来自三处：① 少搬 4/9 的字节；② 少分配/少触碰两个大数组（也少了一半的 TLB 压力与页错误开销）；③ 提高了每条 cache line 上完成的运算量——在带宽受限区间直接转化为加速。
  • 融合的代价：代码可读性下降、失去模块化（讲义特意指出”上面的代码更模块化，比如 Python 的 numpy 数组运算风格”；下面的代码快得多）。这正是”分层的库抽象 vs 单 kernel 手写”的经典权衡，也是为什么现代深度学习编译器要做 operator fusion。
  • 融合不改变算法的工作量（Work），它只改变内存流量；因此它对计算受限的程序几乎无益（Roofline 的平段），只对带宽受限的程序有效（Roofline 的斜段）。

2.7.3 利用共享（exploit sharing）：把访问同一数据的任务放一起

• 定义与目的：把操作同一份数据的任务安排在同一时间、同一个处理器上执行，从而把”多次搬运”变成”一次搬运 + 多次本地使用”。它降低的是固有通信（不是人为通信）。
• 直观解释（”它是什么？”）：类比开小组会。与其让 8 个人各自跑去档案室抄同一份资料（8 趟往返），不如把 8 个人叫到同一间会议室，资料搬一次、大家轮流看。
• 架构/机制图解（CUDA thread block 是硬件层面的”共享”抽象）：

   CUDA 的三层并发 = 三层"共享/通信"粒度
   ┌─────────────────────────── GPU ───────────────────────────┐
   │  SMM/SM core 0            SMM/SM core 1        ...         │
   │  ┌──────────────────┐    ┌──────────────────┐              │
   │  │ Thread block A   │    │ Thread block B   │              │
   │  │  ┌────┐ ┌────┐   │    │  ┌────┐ ┌────┐   │              │
   │  │  │warp│ │warp│   │    │  │warp│ │warp│   │              │
   │  │  └────┘ └────┘   │    │  └────┘ └────┘   │              │
   │  │  __syncthreads() │    │  __syncthreads() │              │
   │  │  shared memory   │    │  shared memory   │              │
   │  │  (块内可见、极快) │    │                  │              │
   │  └──────────────────┘    └──────────────────┘              │
   │        ▲ 同一个 block 的线程**必定**被调度到同一个 SM 上    │
   │        ▲ 因此可以用 shared memory 快速通信/同步（无需走 L2）│
   └────────────────────────────────────────────────────────────┘
   讲义原话：GPU 实现"总是把同一个 block 的线程调度到同一个 GPU core 上"，
   这样块内线程才能利用 shared memory 的低延迟访问与同步。
   编程含义：把"需要共享数据的线程"放进同一个 block，是**免费的通信优化**。

• 性能特征与关键结论：
  • 共享让通信从”跨处理器”降级为”处理器内部/片上”：延迟降一个数量级、带宽升一个数量级，且不增加任何算法工作量。
  • 编译器/运行时的调度策略会直接决定共享是否发生（CUDA block 的硬件映射、OpenMP/线程池的亲和性、操作系统调度）；程序员能做的是把”应该共享的东西”组织进同一个调度单元。

2.7.4 通信粒度（granularity）与人为通信

• 定义与目的：通信的粒度（一次搬多少字节、以及 cache 一致性以什么为单位）会引入人为通信，因此在做数据划分时必须把”最小传输单位”考虑进去。
• 直观解释（”它是什么？”）：类比买调料。你需要 “1 克盐”，但超市只按整包卖（1 包 500 克）——你只能买一整包，多出来的 499 克就是”人为”开销。cache line 就是那个”整包”。
• 架构/机制图解（2D 分块 + cache line 粒度 ⇒ 人为通信；以及伪共享）：

   ① 2D 分块：上下方向的邻居可以整行搬（空间局部性好），
                左右方向的邻居只需要 1 个元素，却必须搬一整条 line：
         ┌───────────────┬───┬───────────────┐
         │  线程 A 的块  │ * │  线程 B 的块  │
         │  (N/√P 宽)    │ ◄ │  (N/√P 宽)    │
         └───────────────┴───┴───────────────┘
                        ▲ 只需要这 1 个元素
         ┌───────────────────────────────────────┐
         │ 但一条 cache line = 4 个元素：        │
         │ [ e ][ e ][ e ][ e ] → 整条都要搬     │  ⇒ 人为通信
         └───────────────────────────────────────┘
         讲义结论：**artifactual communication increases with cache line size**
                   （cache line 越大，左右边界浪费越多）

   ② 伪共享（false sharing）：两边逻辑上完全不共享数据，却因为
      共用一条 cache line 而在硬件层面"共享"了它：
        内存/cache line (64 B) = [ x0 x1 x2 x3 | x4 x5 x6 x7 ] ...
                                  └─ 线程 P1 只写 ─┘└─ 线程 P2 只写 ─┘
         逻辑上：无共享、无竞争（各自写各自的元素，甚至没有数据竞争）
         硬件上：P1 要写 x0 就必须"独占"整条 line，于是 P2 手里的副本失效；
                 P2 要写 x4 又要把 line 抢回独占状态 —— **整条 line 在两个核之间乒乓**
         结果：**固有的通信 = 0，人为的通信 = 每条 line 来回跑**
         （讲义在此处标注："further detail in the upcoming cache coherence lectures"）

   ③ 修法一：blocked data layout（分块布局）消除"跨界"地址
     2D 行主序布局（左）：相邻地址被切到不同处理器 ⇒ 普遍 straddle 分区边界
       P1 P1 P2 P2 P1 P1 P2 P2        ← 连续地址在分区边界反复横跳
       P3 P3 P4 P4 P3 P3 P4 P4
     4D 数组分块布局（右，block-major）：连续地址始终落在同一个分区内
       ┌────────┬────────┐
       │ P1 P1  │ P2 P2  │           ← 每个分区在地址空间里是一段连续区域
       │ P1 P1  │ P2 P2  │
       ├────────┼────────┤
       │ P3 P3  │ P4 P4  │
       │ P3 P3  │ P4 P4  │
       └────────┴────────┘
      注意区分两件事（讲义特意提醒）：
        - blocked assignment of WORK to threads（两种布局都一样）
        - blocked data layout in the ADDRESS SPACE（只有右边是这样）
   ④ 修法二（伪共享专用）：padding，让不同线程写的数据落在不同的 cache line 上

• 性能特征与关键结论：
  • 通信粒度与 cache 一致性粒度都要在”任务划分”阶段考虑；数据布局（layout）与工作划分（assignment）必须联合设计，否则工作划分得再漂亮也会被 cache line 边界吃掉。
  • 伪共享是”代码正确、性能崩塌”的典型：它不产生数据竞争（每个线程只写自己的数据），却产生与”所有线程争抢同一个变量”几乎等价的硬件通信量。

2.8 竞争（contention）：比”通信多少”更致命的是”通信何时发生”

• 定义与目的：竞争指在很短的窗口内对同一个资源发出大量请求的现象，此时该资源成为热点（hot spot）。每个资源都有一个固有吞吐量（单位时间的交易数）：内存、通信链路、服务器、Office hours 的助教都可以这样看。竞争不改变工作量，却让每次操作的完成时间显著变长（排队）。
• 直观解释（”它是什么？”）：讲义用答疑时间作类比。假设”完成一次答疑”由三步组成：① 从咖啡厅走到办公室 5 分钟；② 排队等待（不确定）；③ 教授给出深刻解答 5 分钟。
  • 情形一（有预约）：两位同学分别约了 3:00 与 4:30。两人各走 5 分钟、各被解答 5 分钟，每人耗时都是 10 分钟，谁都不用等。
  • 情形二（3:00–3:20 敞开答疑，不预约）：5 位同学都在 3:00 前后到，只有一个执行资源（一位教授）。第一位同学仍然只花 10 分钟；最后一位同学要排队，他的耗时 = 5 分钟走路 + 排队（几位同学 × 每人 5 分钟）+ 5 分钟答疑，于是变成 23 分钟甚至更久。
  • 结论：竞争让”总操作时间”因为排队而膨胀，而且膨胀量取决于到达的时间分布——这就是为什么”错峰访问（stagger）”是有效手段。
• 架构/机制图解（平坦 vs 树形通信；分布式工作队列）：

   ① 更新一个共享变量：平坦（flat）vs 树形（tree）
      平坦：所有 P 个处理器都直接访问同一个计数器/内存位置
        P0 ─┐
        P1 ─┼──► [ 热点：单个计数器 ]      竞争量 ∝ P（串行化）
        P2 ─┤        吞吐 = 1/t_op         无竞争时延迟最低，但 P 大时排队爆炸
        P3 ─┘
      树形：先两两合并，再逐级上合
        P0 ─┐                       无竞争时延迟更高（log P 级），
        P1 ─┴─►[合并]─┐             但**竞争被分散**：每一级的每个节点
        P2 ─┐         ├─►[合并]─► 总和  只承受常数个请求
        P3 ─┴─►[合并]─┘
        讲义原话：tree structured communication 降低竞争（代价是无竞争时延迟更高）；
                  flat communication 竞争风险高（但无竞争时延迟最低）

   ② 分布式工作队列（distributed work queues）：把"竞争"变成"只在必要时同步"
                 T1        T2        T3        T4        ← worker 线程
               ┌────┐   ┌────┐   ┌────┐   ┌────┐
        子问题→│ Q1 │   │ Q2 │   │ Q3 │   │ Q4 │        ← 每个线程自己的队列
         (task)└─┬──┘   └─┬──┘   └─┬──┘   └─┬──┘
               push/pop         Steal!◄──────┘          ← 本地队列空时才去"偷"
                （无竞争）      （此时线程本来也要闲着，同步的代价可以忽略）
        规则：Pull/Push 只碰**自己的**队列；只有本地队列为空时才从别人的队列尾部/头部偷
        收益：① 有活干时零竞争；② 工作窃取天然负载均衡；③ 偷的工作在"调用树的上层"，
                粒度更大，摊销了窃取成本（详见第 8 讲 continuation stealing）

• 降低竞争的四类手段（讲义最后的总结，也适用于降低通信成本的整体框架）：

2.9 案例研究：把 100 万个粒子放进 16 个格子里（五种解法）

• 问题设定：把 100 万个点粒子按二维位置放进 16 个均匀格子（构建”二维列表的数组”）——这是并行数据结构操作的典型难题（不规则、需要动态写入）。讲义给出的机器是 GTX 980 GPU：每个 SMM core 最多 2048 条 CUDA 线程，GPU 共 16 个 SMM core。这个结构的常见用途是 N-body 问题：给定粒子，找出半径 R 内的所有邻居；把格子边长设成 R，则只需检查周围格子。
• 五种解法的对照（讲义逐条给出了优缺点）：

• 架构/机制图解（解法 ⑤：数据并行分桶的三步）：

   输入:                            步骤 1: 并行算格子号 (per particle)
   particle_index: [ 0  1  2  3  4  5  6  7  8  9 10 11 ... ]   ← 粒子编号
   grid_index:     [ 0  1  2  4  5  3  9  6  6  4  6  4 ... ]   ← 每个粒子落在哪个格子
                                │
                                ▼  步骤 2: 按 grid_index 排序（并行计数排序/基数排序）
   particle_index: [ 0  2  5  2  5  6  9  3  5  1  2  4 ... ]   ← 重新排列后的粒子编号
   grid_index:     [ 0  1  2  2  3  4  4  4  5  6  6  6 ... ]   ← 现在是**有序**的
                                │
                                ▼  步骤 3: 并行求每段的起止（对每个 index 执行一次）
     cell = grid_index[index]
     if (index == 0)                      cell_starts[cell] = index;
     else if (cell != grid_index[index-1]) { cell_starts[cell] = index;
                                             cell_ends[grid_index[index-1]] = index; }
     if (index == numParticles-1)          cell_ends[cell] = index + 1;   // 末端不包含
                                │
                                ▼
     cell_starts: [ 0  1  2  4  5  8  ... ]     每格在排序数组中的起点
     cell_ends  : [ 1  2  4  5  8 11  ... ]     每格在排序数组中的终点（不含）
     ⇒ 要遍历"格子 c 里的所有粒子"：for (i = cell_starts[c]; i < cell_ends[c]; ++i)
                                      访问 particle_index[i]
   代价：一次排序（对比并行、带宽开销）+ 数组的额外存储；
   收益：**完全不需要细粒度同步**，且每个阶段都保持极大并行度（适合 GPU）

• 性能特征与关键结论：
  • 五种解法构成一条清晰的权衡曲线：并行度 ↔ 竞争 ↔ 额外工作/内存。① 牺牲并行度换取零竞争；② 用最大竞争换取零额外工作；③ 用细粒度锁减少竞争（约 16×）；④ 用内存（N 张网格）与合并工作换取”块内廉价同步”；⑤ 用一次排序与额外带宽换取零同步。
  • 没有一种解法普适：在只有 16 个执行单元的机器上，② 也许可以接受（争抢者少）；在有 2048×16 条线程的 GPU 上，只有 ④ 或 ⑤ 能跑满。
  • 这条案例线在后续课程中会继续被展开：第 15/16 讲（同步实现、无锁编程）会解释为什么细粒度锁与无锁结构能降低竞争，第 22~23 讲（并行深度学习）里”用一次排序/直方图替代原子操作“的思路会再次出现（例如 embedding 梯度聚合）。

3. 代码示例与性能分析

本节共 5 个完整示例，分别对应本讲五条主线：① 显式通信的重叠（MPI 非阻塞）、② 算术强度与循环融合、③ 遍历顺序/分块（时间局部性）、④ 伪共享与数据布局、⑤ 竞争（锁的粒度 vs 数据并行）。 除示例 1（需要 MPI 运行时）外，其余示例均用 g++ -O3 -march=native -fopenmp 编译（编译均已验证通过），运行参数写在各自的编译注释里。 给出的实测数字来自一台 双路 AMD EPYC 7V13（64 核/路 ×2，共 128 个线程；L1d 32 KB/核、L2 512 KB/核、L3 32 MB/CCX；AVX2，无 AVX-512；64 B cache line），但这是一台共享机器（同时有其他用户任务，load average 在 8~30 之间波动），绝对时间可能相差数倍。因此文中一律给出多次测量的区间，并强调：只有同一台机器、同一时段、成对对比得到的比值才是可靠结论；绝对值仅供参考。

3.1 示例一：MPI 非阻塞 ghost 行交换（把通信藏到计算背后）

// 编译: mpicxx -O3 -march=native grid_mpi.cpp -o grid_mpi
// 运行: mpirun -np 4 ./grid_mpi 2048 200
#include <mpi.h>
#include <cstdio>
#include <cstdlib>
#include <cmath>
#include <vector>

enum { MSG_ROW = 100 };   // 行消息统一用同一个 tag（发送方与接收方必须匹配）

int main(int argc, char** argv) {
    MPI_Init(&argc, &argv);
    int rank, nprocs;
    MPI_Comm_rank(MPI_COMM_WORLD, &rank);
    MPI_Comm_size(MPI_COMM_WORLD, &nprocs);

    const int N    = (argc > 1) ? atoi(argv[1]) : 2048;
    const int ITER = (argc > 2) ? atoi(argv[2]) : 200;
    const int rows = N / nprocs;            // 本进程拥有的内部行数（假设整除）
    const int ld   = N + 2;                 // 行主序，左右各留 1 列 halo
    const int size = (rows + 2) * ld;       // 本地数组 = 2 行 ghost + rows 行内部

    std::vector<float> A(size, 0.0f);
    const int r0 = rank * rows;             // 本进程负责的全局行区间 [r0, r0+rows)
    for (int i = 1; i <= rows; ++i)
        for (int j = 1; j <= N; ++j) {
            double x = (double)(r0 + i) / N, y = (double)j / N;
            A[i * ld + j] = (float)std::exp(-((x - .5) * (x - .5) + (y - .5) * (y - .5)) / 0.02);
        }

    // 更新一行，返回该行的 |Δ| 之和
    auto update_row = [&](int i) -> float {
        const float* up = &A[(i - 1) * ld];
        const float* dn = &A[(i + 1) * ld];
        float* cur = &A[i * ld];
        float d = 0.0f;
        for (int j = 1; j <= N; ++j) {
            float prev = cur[j];
            float v = 0.2f * (up[j] + dn[j] + cur[j - 1] + cur[j + 1] + cur[j]);
            cur[j] = v;
            d += std::fabs(v - prev);
        }
        return d;
    };

    float mydiff = 0.0f;
    MPI_Barrier(MPI_COMM_WORLD);
    double t0 = MPI_Wtime();

    for (int it = 0; it < ITER; ++it) {
        MPI_Request req[4];
        MPI_Status  st[4];
        int nreq = 0;
        const int up = rank - 1, down = rank + 1;

        // (1) 一次性投递全部非阻塞通信：收 ghost 行 + 发边界行（发的是上一轮的结果）
        if (up >= 0) {
            MPI_Irecv(&A[0 * ld + 1],          N, MPI_FLOAT, up,   MSG_ROW, MPI_COMM_WORLD, &req[nreq++]);
            MPI_Isend(&A[1 * ld + 1],          N, MPI_FLOAT, up,   MSG_ROW, MPI_COMM_WORLD, &req[nreq++]);
        }
        if (down < nprocs) {
            MPI_Irecv(&A[(rows + 1) * ld + 1], N, MPI_FLOAT, down, MSG_ROW, MPI_COMM_WORLD, &req[nreq++]);
            MPI_Isend(&A[rows * ld + 1],       N, MPI_FLOAT, down, MSG_ROW, MPI_COMM_WORLD, &req[nreq++]);
        }

        // (2) 与通信重叠：先算不依赖 ghost 行的内部行 i = 2 .. rows-1
        float d = 0.0f;
        for (int i = 2; i <= rows - 1; ++i) d += update_row(i);

        // (3) 等通信落地（发送缓冲马上就要被复写，必须先 Wait）
        if (nreq) MPI_Waitall(nreq, req, st);

        // (4) 再算边界行：i=1 与 i=rows 依赖 ghost 行
        const int first = 1, last = (rows >= 2) ? rows : 1;
        for (int i = first; i <= last; ++i)
            if (rows < 2 || i == first || i == last) d += update_row(i);

        mydiff += d;
    }

    double local_t = MPI_Wtime() - t0, gtime = 0.0;
    float  gdiff = 0.0f;
    MPI_Allreduce(&mydiff, &gdiff, 1, MPI_FLOAT, MPI_SUM, MPI_COMM_WORLD);
    MPI_Allreduce(&local_t, &gtime, 1, MPI_DOUBLE, MPI_MAX, MPI_COMM_WORLD);
    if (rank == 0)
        printf("P=%d N=%d iters=%d  time=%.3f s  sum|diff|=%.3e  (%.1f Melem-updates/s)\n",
               nprocs, N, ITER, gtime, gdiff, (double)N * N * ITER / gtime / 1e6);
    MPI_Finalize();
    return 0;
}

• 【代码做什么？】
  1. 分解与分配：把 N×N 网格按行做 1D blocked assignment；进程 rank 负责全局行区间 [rank*rows, (rank+1)*rows)。本地数组每行多留 2 列（ld = N+2，左右 halo），上下各多留 1 行（ghost）。
  2. 投递通信：每个进程向上、向下邻居各发一行（自己最上面一行 row 1 给上方、最下面一行 row rows 给下方），同时从两个邻居各收一行到 ghost 行。四个收发动作连续投递，全部是非阻塞的（MPI_Isend/MPI_Irecv 立即返回），用的是同一个 tag MSG_ROW——发送方与接收方的 tag 必须匹配，否则 Waitall 会永久挂起（这是最常见的 MPI 调试陷阱）。
  3. 重叠计算：先算不依赖 ghost 行的内部行 i = 2 .. rows-1。这段时间里消息正在网络上传输，构成 overlap。
  4. 同步点：MPI_Waitall 等所有收发完成。必须在复写发送缓冲之前等待——发送缓冲就是 row 1 与 row rows 这两行，它们马上要在第 (4) 步被更新。这是非阻塞通信的第一号陷阱（”发给 send 的缓冲区在发送完成前不能修改”）。
  5. 算边界行：i=1 与 i=rows 需要 ghost 行数据，放在 Waitall 之后。
  6. 收敛判定：每轮累加本地 \|Δ\|，最后用 MPI_Allreduce 求和（用消息实现全局归约）；真实求解器会把这一步放进循环并按容差提前退出。
• 【并行机制与性能解说】
  • 谁在并行：nprocs 个 MPI 进程，每个进程是独立的地址空间（可能在不同节点上）。数据在进程间只能通过消息流动；rows 行是私有的，ghost 行是复制品。
  • 共享数据如何处理：没有共享内存。每一轮必须做一次边界行复制（ghost 行更新），这就是本讲的固有通信：每个进程每轮收发 2 × N × sizeof(float) 字节（只有两个邻居，不是四个——上下由邻居提供，左右由本地 halo 列解决）。
  • Work / Span / 并行度：
    • 每一步迭代的 Work = Θ(N²) 次单元更新（全网格），每个单元 5 次浮点运算 + 5 次 load + 1 次 store。
    • 每一步迭代的 Span（关键路径） = Θ(N/P)（每个进程内部按行顺序推进，Gauss-Seidel 风格的行间依赖使同一进程的 R 行必须顺序完成）+ L（一次消息的延迟，即 T0 + n/B）。
    • 总 Work = N²·ITER，总 Span = ITER·(N/P + L)。并行度 = Work/Span ≈ N·P（当 N/P ≫ L 时），也就是当每个进程分到的行数足够多时并行度才够。
    • 在这个示例的默认参数下（N=2048, P=4），每进程 rows=512 行，每行 2048 个单元：每轮计算量 ≈ 512×2048 = 1.05 M 次单元更新，而通信量是 2×2048×4 B = 16 KB——字节数之比约为 1.05M×24 B(读写) : 16 KB ≈ 1570 : 1，通信完全不是瓶颈，overlap 甚至不重要。这正是”表面-体积比（surface-to-volume）”的体现：P 变大时 rows 变小，通信/计算比按 1/rows 上升。
  • 瓶颈在哪：
    1. P 很大时通信比例上升：rows = N/P，固有通信量固定为 2N 元素，而计算量 N²/P 随 P 下降 ⇒ 通信占比 ∝ P/N。当 N=2048, P=1024 时 rows=2，通信已经是主导（数值算例见 §4.1）。
    2. 内存带宽：每单元 5 次 load + 1 次 store，算术强度极低（5 flops / 24 B ≈ 0.21 flop/B），单节点内就已经是带宽受限，加进程不会改善单节点内存带宽。
    3. 同步开销：Waitall 是每轮一个硬同步点；如果通信不能被计算完全遮住（比如 rows 太小），进程就会在 Waitall 上排队等待。把 Waitall 放在最晚的位置（先算内部行）就是本示例最重要的优化。
    4. 负载不均：1D blocked 分配下每进程行数相同，负载是均衡的；但如果 N 不能被 nprocs 整除，就会有人多一行。

3.2 示例二：循环融合把”算术强度”从 1/3 提到 3/5

// 编译: g++ -O3 -march=native -fopenmp fusion.cpp -o fusion
// 运行: OMP_NUM_THREADS=8 ./fusion 134217728      (1.28 亿个 float = 512 MB/数组)
#include <omp.h>
#include <cstdio>
#include <cstdlib>
#include <vector>
#include <algorithm>

static void add(int n, const float* A, const float* B, float* C) {
    for (int i = 0; i < n; ++i) C[i] = A[i] + B[i];
}
static void mul(int n, const float* A, const float* B, float* C) {
    for (int i = 0; i < n; ++i) C[i] = A[i] * B[i];
}
static void add_omp(int n, const float* A, const float* B, float* C) {
#pragma omp parallel for schedule(static)
    for (int i = 0; i < n; ++i) C[i] = A[i] + B[i];
}
static void mul_omp(int n, const float* A, const float* B, float* C) {
#pragma omp parallel for schedule(static)
    for (int i = 0; i < n; ++i) C[i] = A[i] * B[i];
}

// ① 三次遍历 + 两个中间数组（"模块化"写法，类似 numpy 的数组运算）
static void three_pass(int n, const float* A, const float* B, const float* C,
                       const float* D, float* E, float* t1, float* t2, int nt) {
    if (nt == 1) { add(n, A, B, t1); mul(n, t1, C, t2); add(n, t2, D, E); }
    else         { add_omp(n, A, B, t1); mul_omp(n, t1, C, t2); add_omp(n, t2, D, E); }
}
// ② 循环融合：每个元素 4 次 load + 1 次 store 完成 3 个浮点运算
static void fused(int n, const float* A, const float* B, const float* C,
                  const float* D, float* E, int nt) {
    if (nt == 1) {
        for (int i = 0; i < n; ++i) E[i] = D[i] + (A[i] + B[i]) * C[i];
    } else {
#pragma omp parallel for schedule(static)
        for (int i = 0; i < n; ++i) E[i] = D[i] + (A[i] + B[i]) * C[i];
    }
}

int main(int argc, char** argv) {
    const int n = (argc > 1) ? atoi(argv[1]) : (16 << 20);
    const size_t sz = (size_t)n;
    std::vector<float> A(sz), B(sz), C(sz), D(sz), E(sz, 0.f), t1(sz), t2(sz);
#pragma omp parallel for schedule(static)
    for (int i = 0; i < n; ++i) { A[i] = i * 1e-7f; B[i] = 1.0f; C[i] = 0.5f; D[i] = 2.0f; }

    const double elt = sizeof(float);
    const double B3 = 9.0 * elt * n;   // 3 遍写法：每元素 9 次 4B 搬运
    const double B1 = 5.0 * elt * n;   // 融合写法：每元素 5 次 4B 搬运
    const int nt = omp_get_max_threads();
    const int REP = 3;
    printf("n = %d 个 float/数组 (%d MB)，OpenMP 线程数 = %d，每项取 %d 次最快\n\n",
           n, (int)(n * elt / (1 << 20)), nt, REP);

    auto bench_pass = [&](const char* name, int threads, double bytes) {
        double best = 1e30;
        for (int r = 0; r < REP; ++r) {
            double t = omp_get_wtime();
            three_pass(n, A.data(), B.data(), C.data(), D.data(), E.data(),
                       t1.data(), t2.data(), threads);
            best = std::min(best, omp_get_wtime() - t);
        }
        printf("%-24s %8.2f ms   %7.2f GB/s (模型)   %6.0f MB\n",
               name, best * 1e3, bytes / best / 1e9, bytes / 1e6);
        return best;
    };
    auto bench_fused = [&](const char* name, int threads, double bytes) {
        double best = 1e30;
        for (int r = 0; r < REP; ++r) {
            double t = omp_get_wtime();
            fused(n, A.data(), B.data(), C.data(), D.data(), E.data(), threads);
            best = std::min(best, omp_get_wtime() - t);
        }
        printf("%-24s %8.2f ms   %7.2f GB/s (模型)   %6.0f MB\n",
               name, best * 1e3, bytes / best / 1e9, bytes / 1e6);
        return best;
    };

    double s3 = bench_pass ("3-pass (1 thread)",   1,  B3);
    double s1 = bench_fused("fused  (1 thread)",   1,  B1);
    double p3 = bench_pass ("3-pass (OpenMP)",     nt, B3);
    double p1 = bench_fused("fused  (OpenMP)",     nt, B1);

    printf("\n实测加速比 融合/三遍: 单线程 %.2fx, 多线程 %.2fx\n", s3 / s1, p3 / p1);
    printf("模型: 3-pass 每元素 9 次搬运 (AI = 3/9 = %.3f math/搬运)；"
           "fused 5 次 (AI = 3/5 = %.3f)\n", 3.0 / 9.0, 3.0 / 5.0);
    printf("按字节: 3-pass %.4f flop/B ；fused %.4f flop/B\n",
           3.0 / (9.0 * elt), 3.0 / (5.0 * elt));
    return 0;
}

• 【代码做什么？】
  1. 分配 5 个”真”数组（A,B,C,D,E）与 2 个临时数组（t1,t2），每个 n = 128M 个 float = 512 MB（工作集远超 L3，保证是内存带宽受限而不是 cache 命中游戏）。
  2. three_pass 用三个独立循环完成 E = D + (A+B)*C，中间结果落在 t1、t2；fused 用一个循环完成同样的计算。
  3. 每个版本先热身、再取 3 次中最快（消除 page fault 与调度噪声的影响）。
  4. 打印模型流量（9 次 vs 5 次 4B 搬运）与由此折算的”模型带宽”，最后打印按字节计的算术强度。
• 【并行机制与性能解说】
  • 并行机制：#pragma omp parallel for schedule(static) 把 n 次独立迭代静态切块分给线程；线程数由 OMP_NUM_THREADS 决定。每个迭代的计算完全独立（无依赖、无同步、无共享写），所以是数据并行（data-parallel）：编译器还会把它向量化（AVX2，8 个 float 一批）。
  • Work / Span / 并行度：
    • Work = 3n 次浮点运算 + 9n（或 5n）次元素搬运。
    • Span = 一个线程处理自己那份的 n/P 次迭代（加上循环启动/结束的开销）：Θ(n/P)；对 parallel for 而言，并行循环的默认末尾屏障也在关键路径上（每次调用都有一个 barrier）。
    • 并行度 = Work / Span = 3n / (n/P) = 3P（P 为线程数）。也就是说这份代码的理论并行度与线程数同阶，而且每次迭代只有 3 个浮点运算 ⇒ 并行度很”浅”，机器越大越难喂饱：一旦内存带宽饱和，再加线程没有任何收益（这正是实测中 8 线程以上收益迅速消失的原因）。

  • 瓶颈与实测（本机实测，512 MB/数组）：

    版本	单线程	8 线程	16 线程	模型搬运量
    3-pass（三次遍历 + 2 临时数组）	195 ~ 288 ms	93 ~ 120 ms	154 ms	9×4 B/元素 = 4.83 GB
    fused（一次遍历）	113 ~ 137 ms	34 ~ 39 ms	98 ms	5×4 B/元素 = 2.68 GB
    比值（3-pass / fused）	1.7 ~ 2.1×	2.7 ~ 3.1×	1.57×	1.8×（= 9/5，理论值）

    （16 线程那一列反而更慢，恰好演示了”带宽受限”的特征：当内存带宽已经被打满、且线程被跨插槽调度、机器又被其他任务占用时，加线程不再有收益。）

  • 为什么实测比 1.8× 还高？ 因为融合除了少搬 44% 的字节，还额外：① 少触碰 2 个 512 MB 的临时数组 ⇒ TLB 压力与缺页开销减半；② 单次遍历对同一批 line 的复用更好 ⇒ 每字节内存请求能产生更多计算（MLP 效率更高）；③ 三次遍历之间有 3 个 barrier，而融合只有 1 个。
  • 算术强度对比（按字节）：3-pass = 3 flop / 36 B = 0.083 flop/B；fused = 3 flop / 20 B = 0.150 flop/B。对照 §4.4 的 Roofline，两者都远远落在带宽受限的斜线段（机器 ridge point 约 38 flop/B）——这解释了为什么”减弱内存流量”能直接变成加速，而”优化算术”在这类代码上几乎无意义。

3.3 示例三：遍历顺序与分块（把”3 读 1 写”降到”1 读 1 写”）

// 编译: g++ -O3 -march=native -fopenmp stencil_tile.cpp -o stencil_tile
// 运行: OMP_NUM_THREADS=8 ./stencil_tile 64 8388608 5 16384 4
//       参数: 行数 H, 列数 W, 迭代次数, 瓦片宽 TW, 瓦片高 TH
// 5 点模板 Jacobi 更新（双缓冲）：每次迭代读三个旧行、写一个新行
#include <omp.h>
#include <cstdio>
#include <cstdlib>
#include <cmath>
#include <vector>
#include <algorithm>

// ① 行主序全扫描：每个线程按行往下走
static void sweep_rowmajor(const float* src, float* dst, int H, int W) {
#pragma omp parallel for schedule(static)
    for (int i = 1; i < H - 1; ++i) {
        const float* up = src + (size_t)(i - 1) * W;
        const float* md = src + (size_t)i * W;
        const float* dn = src + (size_t)(i + 1) * W;
        float* out = dst + (size_t)i * W;
        for (int j = 1; j < W - 1; ++j)
            out[j] = 0.2f * (up[j] + dn[j] + md[j - 1] + md[j + 1] + md[j]);
    }
}

// ② 分块遍历：把列方向切成 TW 宽的瓦片，每个线程按行块处理
static void sweep_tiled(const float* src, float* dst, int H, int W, int TW, int TH) {
#pragma omp parallel for schedule(static)
    for (int ib = 1; ib < H - 1; ib += TH) {
        const int ie = std::min(ib + TH, H - 1);
        for (int jb = 1; jb < W - 1; jb += TW) {
            const int je = std::min(jb + TW, W - 1);
            for (int i = ib; i < ie; ++i) {
                const float* up = src + (size_t)(i - 1) * W;
                const float* md = src + (size_t)i * W;
                const float* dn = src + (size_t)(i + 1) * W;
                float* out = dst + (size_t)i * W;
                for (int j = jb; j < je; ++j)
                    out[j] = 0.2f * (up[j] + dn[j] + md[j - 1] + md[j + 1] + md[j]);
            }
        }
    }
}

int main(int argc, char** argv) {
    const int H = (argc > 1) ? atoi(argv[1]) : 64;
    const int W = (argc > 2) ? atoi(argv[2]) : (8 << 20);
    const int ITERS = (argc > 3) ? atoi(argv[3]) : 5;
    const int TW = (argc > 4) ? atoi(argv[4]) : 16384;
    const int TH = (argc > 5) ? atoi(argv[5]) : 4;
    const size_t n = (size_t)H * W;
    std::vector<float> A(n), B(n, 0.f);
#pragma omp parallel for schedule(static)
    for (int i = 0; i < H; ++i)
        for (int j = 0; j < W; ++j) A[(size_t)i * W + j] = std::sin((double)j * 1e-4) + 0.01f * i;

    printf("线程=%d  网格 %d x %d : 每行 %.0f MB, 两个缓冲区共 %.2f GB;  瓦片 %dx%d (工作集 %.1f KB)\n",
           omp_get_max_threads(), H, W, W * 4.0 / (1 << 20), 2.0 * n * 4 / (1u << 30),
           TW, TH, 3.0 * TW * 4 / 1024);

    double brm = 1e30, btl = 1e30;
    for (int it = 0; it < ITERS; ++it) {
        double t = omp_get_wtime(); sweep_rowmajor(A.data(), B.data(), H, W);
        brm = std::min(brm, omp_get_wtime() - t); std::swap(A, B);
        t = omp_get_wtime();        sweep_tiled(A.data(), B.data(), H, W, TW, TH);
        btl = std::min(btl, omp_get_wtime() - t); std::swap(A, B);
    }
    const double el = (double)(H - 2) * (W - 2);
    printf("  行主序 (模型 3 读 + 1 写 = 16 B/元素): %7.1f ms  %6.2f GB/s\n",
           brm * 1e3, 16.0 * el / brm / 1e9);
    printf("  分块   (模型 1 读 + 1 写 =  8 B/元素): %7.1f ms  %6.2f GB/s   加速比 %.2fx\n",
           btl * 1e3, 8.0 * el / btl / 1e9, brm / btl);
    return 0;
}

• 【代码做什么？】 对一张 64 × 8,388,608 的网格（每行 32 MB，双缓冲共 4 GB）做 5 点模板 Jacobi 更新。
  • sweep_rowmajor：行主序——线程拿到连续的若干行，一行一行往下推，每算一行要读 i-1, i, i+1 三行。
  • sweep_tiled：分块——先按行块 ib（每个线程一个行块），再在列方向以 TW=16384 个元素为瓦片宽度，在瓦片内把 TH=4 行全部算完再跳到下一个瓦片。
  • 两者都取多轮迭代中的最小值，最后按模型搬运量折算带宽。
• 【并行机制与性能解说】
  • 并行机制：外层 #pragma omp parallel for 按行（或行块）静态划分，每个线程独立处理自己那段行，不需要任何同步（Jacobi 双缓冲下每次迭代的读是只读、写是只写，天然无依赖）。内层循环被 GCC 自动向量化（AVX2，32 B/次）。
  • Work / Span / 并行度：Work = (H-2)(W-2) 次单元更新（约 5.2 亿次）；Span = 一个线程处理 (H-2)/P 行的 Θ((H-2)/P · W)；并行度 = Work/Span = P（只有 P，没有额外的维度可挖）。迭代之间因为双缓冲而必须顺序执行，所以总 Span = ITERS × (H/P) × W。
  • 重用距离分析（为什么分块有用）：
    • 行主序：输出行 i 需要旧行 i-1, i, i+1。旧行 i 会在输出 i-1、i、i+1 时被用到三次，两次使用之间隔了整整一行（32 MB 读 + 32 MB 写 = 64 MB 的间隔数据量）。本机每核可用的 L3 slice 只有 32 MB ⇒ 重用距离 > cache 容量 ⇒ 三次读全部落回 DRAM，模型搬运量 = 3 读 + 1 写 = 16 B/元素。
    • 分块：瓦片宽 16 K 列 = 64 KB，一个瓦片要用的三行窗口 = 3 × 64 KB = 192 KB，远小于 512 KB 的 L2 ⇒ 每行数据从 DRAM 只取一次、后续两次都在 cache 里命中，模型搬运量 = 1 读 + 1 写 = 8 B/元素。
    • 所以理论加速比上限 = 16/8 = 2×。
  • 实测（8 线程）：行主序 116 ~ 151 ms vs 分块 73 ~ 139 ms，成对测得的加速比在 1.0× ~ 2.1× 之间（一次较干净的测量得到 122 ms vs 65 ms = 1.89×，与 2× 的模型吻合）。必须诚实说明：共享机器上这个比值会掉到 1.0× 附近；瓦片尺寸也很敏感（TW=4096 时只有 1.43×，TW=65536 时 1.59×）。原因是：现代 CPU 的硬件预取器能同时跟踪多个顺序流，行主序的三个读流本身就被预取得很好；只有当程序真正处于带宽受限、且重用距离确实超过最后一级 cache 时，分块才稳定获益。在 GPU 上（shared memory 只有几十 KB、没有大 L3、预取器弱）分块的收益要大得多，这也解释了为什么 GPU 编程几乎总是在做显式分块。
  • 瓶颈与注意事项：① 分块引入额外的循环与边界判断开销（std::min、非整除尾块）；② 瓦片参数必须按每核可用 cache（本例中是 L2 512 KB）调，而不是按整个芯片的 L3 总量调——并行程序里每个线程只能用到自己那份 cache；③ 分块改变了更新的遍历顺序，如果算法本身有顺序依赖（如就地 Gauss-Seidel），数值结果会改变，这不总是可以接受的。

3.4 示例四：伪共享——”逻辑上无竞争”却慢 13 倍

// 编译: g++ -O3 -march=native -fopenmp falseshare.cpp -o falseshare
// 运行: OMP_NUM_THREADS=8 ./falseshare 5000000
// 三个版本工作量完全相同：每线程做 ITER 次"读-改-写"，数据完全私有、逻辑上无竞争
#include <omp.h>
#include <cstdio>
#include <cstdlib>
#include <vector>

struct alignas(64) Padded { volatile long v; };   // sizeof == 64，独占一条 cache line

int main(int argc, char** argv) {
    const int T = omp_get_max_threads();
    const long ITER = (argc > 1) ? atol(argv[1]) : 5000000L;
    omp_set_num_threads(T);

    std::vector<long> bad(T, 0);        // ① T 个 long 紧密排列：8 线程共用 1 条 cache line
    volatile long* badp = bad.data();   //    用 volatile 指针防止编译器把循环折叠掉
    std::vector<Padded> good(T);        // ② 每线程一条独立 cache line
    std::vector<long> ro(1024);         // ③ 只读表：给本地累加制造"真实工作量"
    for (int i = 0; i < 1024; ++i) ro[i] = i & 3;
    volatile long* rop = ro.data();

    double t = omp_get_wtime();
#pragma omp parallel
    {
        const int tid = omp_get_thread_num();
        for (long i = 0; i < ITER; ++i) badp[tid] = badp[tid] + 1;   // 只碰自己的元素
    }
    const double tb = omp_get_wtime() - t;

    t = omp_get_wtime();
#pragma omp parallel
    {
        const int tid = omp_get_thread_num();
        for (long i = 0; i < ITER; ++i) good[tid].v = good[tid].v + 1;
    }
    const double tg = omp_get_wtime() - t;

    long total = 0;
    t = omp_get_wtime();
#pragma omp parallel reduction(+ : total)
    {
        const int tid = omp_get_thread_num();
        long local = 0;
        for (long i = 0; i < ITER; ++i) local += rop[(i + tid) & 1023];   // 纯本地累加
        total += local;
    }
    const double tl = omp_get_wtime() - t;

    printf("线程数 %d，每线程 %ld 次更新（共 %.1f M 次）；三种写法逻辑上都没有数据竞争\n",
           T, ITER, T * ITER / 1e6);
    printf("  ① 相邻计数器（伪共享） : %8.1f ms   %7.1f M 次/秒\n", tb * 1e3, T * ITER / tb / 1e6);
    printf("  ② 每计数器独占一条 cache line : %8.1f ms   %7.1f M 次/秒   加速比 %.1fx\n",
           tg * 1e3, T * ITER / tg / 1e6, tb / tg);
    printf("  ③ 线程本地累加后汇总   : %8.1f ms   %7.1f M 次/秒   加速比 %.1fx  (和 %ld)\n",
           tl * 1e3, T * ITER / tl / 1e6, tb / tl, total);
    return 0;
}

• 【代码做什么？】 三个版本做的工作在逻辑上完全一样：每个线程把自己的计数器自增 ITER 次。
  • ① bad[tid]：T 个 long 在内存里紧挨着（8 线程共 64 B = 正好一条 cache line）。每个线程只写属于自己的那一个元素——没有任何数据竞争。
  • ② Padded：用 alignas(64) 让每个计数器独占一条 64 B cache line（sizeof(Padded) == 64）。
  • ③ 线程本地累加：只用寄存器/栈，最后做一次归约。
• 【并行机制与性能解说】
  • 并行机制：#pragma omp parallel 启动固定线程数，线程 tid 只访问下标 tid 的数据。三种写法在语义上都没有共享写；差别只在内存地址是否落在同一条 cache line 上。
  • 为什么①会慢：缓存一致性以 cache line 为最小单位维护所有权。线程 A 写 bad[0] 需要独占整条 line，于是线程 B 手里的同一 line 副本失效；B 写 bad[1] 又要把 line 抢回独占状态 ⇒ 一条 line 在 8 个核之间来回”乒乓”。这就是伪共享（false sharing）：代码里没有共享，硬件上共享了。
  • Work / Span / 并行度：Work = T × ITER 次自增（常量）；Span = ITER 次自增（每个线程各自跑完自己那份）；并行度 = Work/Span = T——理想情况下应该完美线性加速。伪共享不改变 Work 与 Span（算法层面），它改变的是每次自增的硬件代价：从”L1 命中 + 寄存器累加”（约 0.5 ns）变成”一次 cache line 所有权迁移”（几十到几百 ns）。

  • 实测：

    版本	8 线程	相对①	4 线程	相对①
    ① 相邻计数器（伪共享）	43 ~ 91 ms（≈9 ~ 18 ns/次）	1.0×	22.1 ms	1.0×
    ② 每计数器独占 cache line	2.5 ~ 3.2 ms（≈0.5 ~ 0.64 ns/次）	13 ~ 33×	2.8 ms	7.9×
    ③ 线程本地累加 + 归约	2.4 ms（稳定）	18 ~ 38×	2.4 ms	9.2×

    （② 与 ③ 的绝对时间几乎不随机器负载变化——因为它们根本不产生核间通信；而 ① 在同一台机器被抢得很厉害时曾测得 1097 ms，比值掉到 3.7×。“谁对负载敏感”本身就是伪共享的指纹。）

  • 瓶颈与结论：失速的根源不是锁、不是原子操作、也不是负载不均，而是一致性协议在 cache line 粒度的所有权迁移（第 11~13 讲会给出 MSI/MESI 状态机与目录协议的细节）。修法有三层：① padding（把不同线程写的数据隔离到不同 line）；② 本地复制/批量更新（先在自己的私有累加器上攒够一批再写回共享位置，把 N 次 line 迁移压成 N/k 次）；③ 改数据结构布局（blocked data layout），让连续地址属于同一个线程——这正是讲义”reducing artifactual comm: blocked data layout”一页的内容。
  • 代价与权衡：padding 会浪费内存与带宽（8 个计数器占 512 B 而不是 64 B）；如果计数器很多、又很稀疏地被访问，padding 可能反而增加 cache 占用。实践中更常见的做法是每个线程一个 padded 的局部结构体（alignas(64) 的 struct PerThread { ... }），或者用 OpenMP 的 reduction 子句让编译器/运行时替你处理。

3.5 示例五：竞争——全局锁、每桶锁与数据并行排序

// 编译: g++ -O3 -march=native -fopenmp bins.cpp -o bins
// 运行: OMP_NUM_THREADS=16 ./bins 4194304
// 把 N 个粒子按二维位置放进 16x16=256 个桶（"二维列表的数组"），三种并行策略对比：
//   A 单一全局锁     B 每桶一把锁（细粒度锁）     C 数据并行：计数排序 + 前缀和 + 并行求桶边界
#include <omp.h>
#include <cstdio>
#include <cstdlib>
#include <vector>
#include <algorithm>

static const int G = 16;              // 16x16 网格
static const int NC = G * G;          // 256 个桶

static inline int cell_of(float x, float y) {
    int cx = (int)(x * G); if (cx < 0) cx = 0; if (cx > G - 1) cx = G - 1;
    int cy = (int)(y * G); if (cy < 0) cy = 0; if (cy > G - 1) cy = G - 1;
    return cy * G + cx;
}

int main(int argc, char** argv) {
    const int n = (argc > 1) ? atoi(argv[1]) : (4 << 20);
    std::vector<float> px(n), py(n);
    unsigned s = 12345u;
    for (int i = 0; i < n; ++i) {          // 确定性伪随机位置
        s = s * 1664525u + 1013904223u; px[i] = (s >> 8) * (1.0f / (1 << 24));
        s = s * 1664525u + 1013904223u; py[i] = (s >> 8) * (1.0f / (1 << 24));
    }
    std::vector<int> ref(NC, 0);           // 串行参照结果
    for (int i = 0; i < n; ++i) ref[cell_of(px[i], py[i])]++;

    std::vector<int> next(n), head(NC), cnt(NC);
    printf("粒子数 %d，桶数 %d，OpenMP 线程数 %d\n\n", n, NC, omp_get_max_threads());

    // ---------- A: 单一全局锁（对应讲义"解法 2"）----------
    double tA;
    {
        std::fill(head.begin(), head.end(), -1);
        omp_lock_t gl; omp_init_lock(&gl);
        double t = omp_get_wtime();
#pragma omp parallel for schedule(static)
        for (int i = 0; i < n; ++i) {
            const int c = cell_of(px[i], py[i]);
            omp_set_lock(&gl);
            next[i] = head[c]; head[c] = i;      // 头插法把粒子挂进桶 c 的链表
            omp_unset_lock(&gl);
        }
        tA = omp_get_wtime() - t;
        omp_destroy_lock(&gl);
        std::fill(cnt.begin(), cnt.end(), 0);
        for (int c = 0; c < NC; ++c) for (int p = head[c]; p != -1; p = next[p]) cnt[c]++;
        printf("A 单一全局锁         : %8.1f ms   %s\n", tA * 1e3, (cnt == ref ? "结果正确" : "结果错误"));
    }

    // ---------- B: 每桶一把锁（对应讲义"解法 3"）----------
    double tB;
    {
        std::fill(head.begin(), head.end(), -1);
        std::vector<omp_lock_t> lk(NC);
        for (int c = 0; c < NC; ++c) omp_init_lock(&lk[c]);
        double t = omp_get_wtime();
#pragma omp parallel for schedule(static)
        for (int i = 0; i < n; ++i) {
            const int c = cell_of(px[i], py[i]);
            omp_set_lock(&lk[c]);
            next[i] = head[c]; head[c] = i;
            omp_unset_lock(&lk[c]);
        }
        tB = omp_get_wtime() - t;
        for (int c = 0; c < NC; ++c) omp_destroy_lock(&lk[c]);
        std::fill(cnt.begin(), cnt.end(), 0);
        for (int c = 0; c < NC; ++c) for (int p = head[c]; p != -1; p = next[p]) cnt[c]++;
        printf("B 每桶一把锁(细粒度) : %8.1f ms   %s   (相对 A 加速 %.1fx)\n",
               tB * 1e3, (cnt == ref ? "结果正确" : "结果错误"), tA / tB);
    }

    // ---------- C: 数据并行（对应讲义"解法 5"）：无锁、无原子操作 ----------
    double tC;
    {
        const int T = omp_get_max_threads();
        std::vector<int> hist((size_t)T * NC, 0), off((size_t)T * NC, 0), sorted(n), gidx(n);
        std::vector<int> cs(NC, -1), ce(NC, -1);
        double t = omp_get_wtime();
        // 步骤 1: 每个线程统计自己区间内的私有直方图（无同步）
#pragma omp parallel
        {
            const int tid = omp_get_thread_num(), nt = omp_get_num_threads();
            const int lo = (int)((long)n * tid / nt), hi = (int)((long)n * (tid + 1) / nt);
            int* h = &hist[(size_t)tid * NC];
            for (int i = lo; i < hi; ++i) h[cell_of(px[i], py[i])]++;
        }
        // 步骤 2a: 对"所有桶的总数"做一次排他前缀和 —— 每个桶在输出数组中的全局起点
        std::vector<int> gstart(NC, 0);
        {
            int acc = 0;
            for (int c = 0; c < NC; ++c) {
                gstart[c] = acc;
                for (int t2 = 0; t2 < T; ++t2) acc += hist[(size_t)t2 * NC + c];
            }
        }
        // 步骤 2b: 每个线程再在桶内部叠加"自己前面那些线程"的计数（工作量 T*NC，可忽略）
#pragma omp parallel for schedule(static)
        for (int c = 0; c < NC; ++c) {
            int acc = gstart[c];
            for (int t2 = 0; t2 < T; ++t2) { off[(size_t)t2 * NC + c] = acc; acc += hist[(size_t)t2 * NC + c]; }
        }
        // 步骤 3: 每个线程把自己的粒子散射到自己的私有输出区间（无原子操作、无锁）
#pragma omp parallel
        {
            const int tid = omp_get_thread_num(), nt = omp_get_num_threads();
            const int lo = (int)((long)n * tid / nt), hi = (int)((long)n * (tid + 1) / nt);
            int* myoff = &off[(size_t)tid * NC];
            for (int i = lo; i < hi; ++i) {
                const int c = cell_of(px[i], py[i]);
                const int pos = myoff[c]++;
                sorted[pos] = i; gidx[pos] = c;        // 排序后的粒子表 + 每个位置的桶号
            }
        }
        // 步骤 4: 并行找出每个桶的起止位置（讲义中的 cell_starts / cell_ends 写法）
#pragma omp parallel for schedule(static)
        for (int idx = 0; idx < n; ++idx) {
            const int c = gidx[idx];
            if (idx == 0) {
                cs[c] = idx;
            } else if (c != gidx[idx - 1]) {
                cs[c] = idx;
                ce[gidx[idx - 1]] = idx;               // 末端不包含
            }
            if (idx == n - 1) ce[c] = idx + 1;
        }
        tC = omp_get_wtime() - t;
        std::fill(cnt.begin(), cnt.end(), 0);
        int bad = 0;
        for (int c = 0; c < NC; ++c) {
            cnt[c] = (cs[c] < 0) ? 0 : ce[c] - cs[c];
            if (cnt[c] != ref[c]) bad++;
        }
        printf("C 数据并行(排序)     : %8.1f ms   %s   桶边界不匹配数 = %d   (相对 A 加速 %.1fx)\n",
               tC * 1e3, (bad == 0 ? "结果正确" : "结果错误"), bad, tA / tC);
    }
    return 0;
}

• 【代码做什么？】 用三种策略把 400 万个随机粒子分到 256 个桶里，并与串行参照结果对比正确性。
  • A：所有线程共用一个 omp_lock_t，插入前加锁、插入后解锁（讲义解法 2：全局锁 ⇒ 巨大竞争）。
  • B：每个桶一把锁，只锁自己那个桶（讲义解法 3：细粒度锁 ⇒ 竞争降低约 16 倍，因为 256 个桶 / 16 线程）。
  • C：完全不用锁/原子操作（讲义解法 5）：每个线程先建私有直方图 → 前缀和算出每个桶的全局起点、以及”桶内第 t 个线程的起点” → 每个线程把粒子散射到自己的私有输出区间 → 最后并行地对有序的 gidx 数组求每桶的 cell_starts/cell_ends。
• 【并行机制与性能解说】
  • 并行机制：A/B 用 parallel for 划分粒子（每个粒子一个迭代），同步放在每次插入上；C 划分粒子区间（每个线程一段连续粒子），同步只发生在两处隐式屏障（步骤 1 之后、步骤 3 之后），每次插入都不需要同步。
  • Work / Span / 并行度：
    • A/B：Work = n 次插入（每次 = 计算格子号 + 哈希/链表操作 + 一次串行化访问）；Span = n × t_crit（所有插入都被全局串行化，t_crit 是一次临界区时间）；并行度 = Work/Span ≈ 1——加锁版本在算法层面就没有并行度。B 的 Span 降到 n/NC × t_crit（每个桶内部仍然串行）；但真正的瓶颈往往不是”桶内部的串行”，而是 head[c] 与 next[i] 这些随机地址所在 cache line 的乒乓——这与伪共享是同一个硬件机制。
    • C：Work = n（直方图）+ n（散射）+ n（求桶边界）≈ 3n；Span = n/P（三个并行区域各自）+ O(log P) 的屏障；并行度 ≈ 3P。代价是额外的内存流量：读 px/py（8n 字节）、写 sorted/gidx（8n 字节）、再读一遍 gidx 求桶边界（4n 字节）≈ 20n 字节的额外带宽。

  • 实测（400 万粒子、256 桶）：

    策略	1 线程	4 线程	16 线程（多次测量区间）	16 线程时每次插入的代价
    A 单一全局锁	38.7 ms	341.5 ms	1748 ~ 2530 ms	≈ 437 ~ 633 ns（竞争导致负扩展）
    B 每桶一把锁	57.5 ms	190.0 ms	1109 ~ 1176 ms	≈ 277 ~ 294 ns（仍然负扩展）
    C 数据并行（排序）	33.7 ms	16.6 ms	4.5 ~ 6.9 ms	≈ 1.1 ~ 1.7 ns（接近线性扩展）

  • 关键解读：
    1. 加锁版本随线程数增加而变慢（A：1 线程 38.7 ms → 16 线程 1748 ~ 2530 ms，慢 45 ~ 65 倍）。这不是”没有加速”，而是负扩展：竞争 + cache line 迁移把每次插入的代价从 1 线程时的 ~10 ns 抬到 633 ns。“锁的粒子（granularity）”只把竞争降低 2.2 倍，而不是理想的 16 倍，因为每个粒子随机落到某个桶，桶头指针所在的 line 仍然在不同核之间反复迁移。
    2. 数据并行版本的实测时间（4.5 ~ 6.9 ms）几乎正好等于它额外需要的带宽代价：20n = 80 MB，折算出的等效带宽是 80 MB / 4.5 ms ≈ 17.8 GB/s（负载高时为 80 MB / 6.9 ms ≈ 11.6 GB/s），正落在这台机器流式带宽的量级上。也就是说：C 用”一次排序的带宽”买到了”零同步”——这正是讲义对解法 5 的评价（”maintains a large amount of parallelism and removes the need for fine-grained synchronization… at cost of a sort and extra passes over the data”）。
    3. 讲义对解法 ④（部分结果 + 合并）的定位也在这里得到印证：把同步”降级”到片上/块内（CUDA shared memory）比”减少锁的粒度”更有效，因为片上同步的绝对成本低了一个数量级。
  • 瓶颈总结：A/B 的瓶颈是竞争 + 一致性通信（不是工作量）；C 的瓶颈是额外带宽（不是同步）。选择哪种方案，取决于”竞争成本”与”带宽成本”哪个更贵——而这又取决于机器（GPU 上带宽更贵、锁更贵，所以更倾向 ④/⑤；16 核 CPU 上两者都可能可行）。

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

本节把前面的定性讨论变成可计算的模型。所有算例都显式给出假设参数，便于替换成你机器上的真实数字。

4.1 通信代价模型：T(n) = T0 + n/B 与有效带宽

• 公式：T(n) = T0 + n/B；有效带宽 BW_eff(n) = n / T(n) = B / (1 + T0·B/n)。
• 达到带宽一半所需的最小消息：n ≥ T0·B。
• 数值算例 A（1D 分块网格求解器的每轮通信）：
  • 假设：N = 4096，P = 64 个进程（1D blocked ⇒ 每进程 rows = N/P = 64 行）；每个进程每轮收发 2 行 × 4096 × 4 B = 32 KB；网络启动延迟 T0 = 2 μs，链路带宽 B = 6.4 GB/s。
  • T = 2 μs + 32768 / 6.4e9 s = 2 μs + 5.12 μs = 7.12 μs；有效带宽 = 32 KB / 7.12 μs = 4.60 GB/s，仅为 B 的 72%（另 28% 的时间花在启动上）。
  • 若把消息拆成 64 条、每行一个元素：T_each = 2 μs + 4/6.4e9 ≈ 2.0006 μs，总时间 = 64 × 2.0006 μs × 2（收+发）≈ 256 μs——比合并成一条慢 36 倍。这就是”send fewer messages, make messages larger“的定量依据。
  • 计算侧：每进程每轮 64 × 4096 = 262,144 次单元更新，按每次 2 个周期（SIMD、带宽受限）@3 GHz 计 ⇒ 174.8 μs。通信/计算 = 7.12/174.8 = 4.1%，即使完全不重叠也只损失 4%。
  • 规模恶化：把 P 提到 512（rows = 8）⇒ 计算降到 8 × 4096 × 2/3e9 = 21.8 μs，而通信仍是 7.12 μs ⇒ 33%；解得 rows ≈ 2.6（即 P ≈ 1575）时通信与计算相等——强扩展（fixed N）下必然存在一个无法逾越的 P 上限。

4.2 Work-Span 模型与贪婪调度

• 公式：T_P ≤ W/P + S（对任意贪婪调度器成立）；并行度 = W / S；效率 = (W/P) / T_P ≥ 1/(1 + S·P/W)。
• 判据：只要 P ≪ W/S，程序就是并行度充足的，瓶颈不会在关键路径上；一旦 P 接近 W/S，加核就不再有效。
• 数值算例 B（把”通信延迟”翻译成”等效工作量”）：
  • 沿用 §4.1 的参数，把一次消息延迟折算成等效的单元更新数：每单元更新 2 cycles / 3 GHz = 0.667 ns，故 L = 7.12 μs ≈ 10,680 次单元更新。
  • MPI 网格求解器（N=2048，ITER 轮）的 W = N²·ITER，”关键路径”里串行的是每轮的 N/P 行加上一次消息延迟：S = ITER·(N/P + L)。
  • 效率 ≈ 1 / (1 + (N + P·L)/N²)：
    • P = 64, N = 2048：P·L/N² = 64 × 10680 / 4.19e6 = 0.163 ⇒ 效率 ≈ 86%。
    • P = 1024, N = 2048：P·L/N² = 1024 × 10680 / 4.19e6 = 2.61 ⇒ 效率 ≈ 28%。
  • 结论：“延迟”在小规模时被计算量淹没，在大规模时变成主项。这也解释了为什么超算上要弱扩展（weak scaling，保持每进程的问题规模不变），而不是强扩展。
• 对比示例二的 Work/Span：parallel for 里每次迭代独立，所以 S = Θ(1)（一次迭代的工作量，加上一次屏障），W = 3n ⇒ 并行度 = Θ(n) ≈ 1.34 亿，远超任何机器的核数。这类代码的瓶颈绝不是 Span，而是带宽：由 §4.4 的带宽模型算出的下界是 4.83 GB / 20 GB/s = 241 ms（3-pass）与 2.68 GB / 20 GB/s = 134 ms（fused），而实测为 195 ~ 288 ms 与 113 ~ 137 ms——与模型下界同一量级（部分测量甚至比 20 GB/s 的假设更快，说明这台机器的实际流式带宽高于 20 GB/s），说明实测到的差别几乎完全由”少搬字节”解释。

4.3 Little’s law 与 memory-level parallelism（MLP）

• 公式：带宽 = 并发请求数 / 延迟，即 B = MLP / L，或 MLP = B·L / 每次请求的字节数。（CS149 版本强调的结论：稳态下处理器的利用率只取决于指令吞吐与内存吞吐，与内存延迟、与”在途请求数”本身无关——前提是有足够的在途请求把带宽填满。）
• 数值算例 C：
  • 假设要打满 B = 20 GB/s，cache line = 64 B ⇒ 需要 312.5 M 次 cache line 传输/秒。
  • 若 DRAM 延迟 L = 200 cycles @ 3 GHz = 66.7 ns：在途 miss 数 = 312.5e6 × 66.7e-9 ≈ 20.8。
  • 若访问的是远端 NUMA 内存（L = 400 cycles = 133 ns）：在途 miss 数 ≈ 41.7。
  • 把这 21~42 个在途请求分摊到 16 个核上，每核只需 1.3 ~ 2.6 个未完成的 load——很容易做到；这就是为什么现代 CPU 靠乱序执行 + 预取可以打满内存带宽。
  • 反过来：如果一个核只有 1 个在途 miss（严格顺序代码），它能达到的带宽是 64 B / 66.7 ns = 0.96 GB/s——只有机器带宽的 5%。这就是”计算没问题、就是慢 20 倍”的根源（也是为什么”把数组访问全改成 A[0]“这个 high watermark 实验能立刻告诉你局部性还有多少油水）。
  • 顺带算一个重要数字：一次遍历 256 MB 数组，在 20 GB/s 下至少要 12.8 ms；如果是”读 + 写”（如 E[i] = f(A[i])），流量翻倍 ⇒ 25.6 ms。这就是任何”每轮全量扫描”算法的硬下界。

4.4 Roofline 模型与算术强度

• 公式：可达性能 ≤ min(峰值算力, 算术强度 × 峰值带宽)；两条线的交点称为 ridge point：I_ridge = 峰值算力 / 峰值带宽。
• 架构/机制图解（ASCII Roofline）：

   性能 (GFLOP/s)
    ▲
768 ┤━━━━━━━━━━━━━━━━━━━━━┳━━━━━━━━━━━━━━━━━━━━━━━━  ← 峰值算力 (水平段: compute limited)
    │                     ┃
    │   斜线段             ┃   水平段
    │  (memory BW limited) ┃
600 ┤                     ┃
    │                  ╱  ┃
400 ┤               ╱     ┃
    │            ╱        ┃
200 ┤         ╱           ┃
    │      ╱              ┃
  ┤   ╱                   ┃
  0 └──┬────┬────┬────┬───┺────┬────┬────┬────┬──────► 算术强度 (flop/byte)
       0.15 0.31 0.63 1.0  38.4  64    128   256
        │    │    │    │     ▲
        │    │    │    │     └── ridge point = 768 / 20 = 38.4 flop/B
        │    │    │    └── 分块矩阵乘 (b=256): AI = 64 → 计算受限
        │    │    └── 分块模板 (8 B/元素, 5 flops): AI = 0.625
        │    └── 行主序模板 (16 B/元素): AI = 0.3125
        └── 融合逐元素运算 (20 B/元素, 3 flops): AI = 0.15
            三遍逐元素运算 (36 B/元素): AI = 0.083（更低，未画出）

   说明：斜线的斜率就是峰值带宽 B；"提高算术强度"= 在图上向右移动，
         一旦越过 ridge point 就进入水平段（此时再减少内存流量没有收益）。

• 数值算例 D（在这台机器上把每条曲线算出来）：
  • 峰值算力：16 核 × 3.0 GHz × 8 (AVX 浮点宽度) × 2 (FMA) = 768 GFLOP/s （注：这里按”每 lane 每周期 2 flop 的 FMA”计一次；若按两个 FMA 端口计则为 1536 GFLOP/s。硬件参数换成 P 核 × f × 宽度 × 2 即可。）
  • 峰值带宽：B = 20 GB/s ⇒ ridge point = 768/20 = 38.4 flop/B。
  • 融合逐元素运算（示例二）：3 flop / 20 B = 0.15 flop/B ⇒ 上界 0.15 × 20 = 3.0 GFLOP/s，只有峰值的 0.39%。
  • 三遍逐元素运算：3/36 = 0.083 flop/B ⇒ 1.67 GFLOP/s。
  • 模板计算：每元素 5 个浮点运算；行主序 16 B/元素 ⇒ 0.3125 flop/B → 6.25 GFLOP/s；分块 8 B/元素 ⇒ 0.625 flop/B → 12.5 GFLOP/s（这正是示例三中 2× 加速的来源）。
  • 分块矩阵乘：2n³ 次浮点运算，DRAM 流量 ≈ 8n³/b 字节（b = 块边长，4 B/元素）⇒ AI = b/4 flop/B。要越过 ridge point 需 b ≥ 4 × 38.4 = 154：b = 64（进 L1）时 AI = 16 ⇒ 仍受带宽限制（上界 320 GFLOP/s）；b = 256（进 L2/L3）时 AI = 64 ⇒ 进入计算受限区。
  • 结论：本讲的四条优化（融合、分块、共享数据、增大消息粒度）在 Roofline 上的作用都是把工作点向右推；只有越过了 ridge point，才算”真的解决了局部性问题”。

4.5 竞争与排队的定量：锁的串行化

• 模型：设一次临界区（含 cache line 所有权迁移）需要 t_crit，则单一全局锁的吞吐上限为 1/t_crit，与线程数无关；线程数越多，等待时间越长（T ≈ 到达顺序 × t_crit），于是出现负扩展。
• 数值算例 E（用 §3.5 的实测数据反推成本）：
  • 全局锁：4 M 次插入 / 1748 ~ 2530 ms ⇒ 每次插入 437 ~ 633 ns，吞吐 1.58 ~ 2.29 M 次/秒（1 线程时是 38.7 ms ⇒ 每次 9.7 ns）。45 ~ 65 倍的膨胀全部来自竞争与一致性迁移（不是算法工作量）。
  • 每桶锁（256 把锁、16 线程，平均每桶 1 个线程）：1109 ~ 1176 ms ⇒ 每次 277 ~ 294 ns，吞吐 3.4 ~ 3.6 M 次/秒 ⇒ 只比全局锁好 1.6 ~ 2.2 倍，远低于”竞争降低 16 倍”的理想值——因为桶头指针 head[c] 所在的 cache line 仍在核间迁移，而锁的原子操作本身也不便宜。
  • 数据并行：4.5 ~ 6.9 ms ⇒ 每次插入 1.1 ~ 1.7 ns。它的成本不在同步，而在额外带宽：20n = 80 MB（读 px/py 32 MB + 写 sorted/gidx 32 MB + 读 gidx 16 MB），80 MB / 4.5 ms ≈ 17.8 GB/s（负载高时 ≈ 11.6 GB/s）——正落在这台机器流式带宽的量级上，说明它已经被带宽打满。
  • 决策规则：把两种方案的成本写成同一单位（时间）再比较——加锁方案的成本 ≈ 每次操作 × 竞争系数 t_crit；数据并行方案的成本 ≈ 额外字节数 / 带宽。本例中 1.1 ns 对 289 ns，差约 260 倍，数据并行完胜；但如果桶数只有 2 个、线程数也只有 2 个，竞争系数很小，额外的排序带宽反而会成为纯亏损。

4.6 Amdahl 定律与”固定问题规模”的陷阱

• 公式（Amdahl）：设程序中无法并行/被串行化的比例为 f，则 Speedup(P) ≤ 1 / (f + (1-f)/P)，上界 1/f。
• 数值算例 F-1：
  • f = 5%：P = 16 时 1/(0.05 + 0.95/16) = 9.1×；无论 P 多大，上限只有 20×。
  • f = 1%：P = 64 时 1/(0.01 + 0.99/64) = 39.4×（上限 100×）。
  • 在并行程序里，”1% 的串行”经常就是”通信/同步的那 1%”：所以本讲的通信优化同时也在压低 f。
• 数值算例 F-2（固定问题规模的扩展陷阱）：
  • 1D 分块的每处理器通信/计算元素比 = 2N / (N²/P) = 2P/N：
    • N = 258, P = 32：2 × 32 / 258 = 0.248 ⇒ 每算 4 个元素就要搬 1 个元素（还要加上 halo 行带来的 25% 冗余计算：rows = 8，其中 2 行是边界），通信与同步开销完全吃掉并行收益——这正是 CS149 版本给出的”SGI Origin 2000 上 258×258 网格在 32 处理器上几乎没有加速”的成因。
    • N = 1024, P = 32：2 × 32 / 1024 = 0.0625 ⇒ 每 16 个计算元素才搬 1 个，扩展性良好（同一台机器上 1K×1K 网格扩展良好）。
  • 推论：评估并行机时必须同时报告问题规模；小问题在大型并行机上”没有加速”并不说明机器差，只说明问题的算术强度相对于机器规模太低。
• 超线性加速（super-linear speedup）与”伪加速”：
  • 真超线性：问题的工作集在单处理器上装不进 cache、在多处理器上却能分摊到每个处理器的私有 cache里（每个处理器只要算自己那一小块），于是单机跑的是”cache miss 版”、并行跑的是”cache hit 版”。
  • 假超线性：单机上工作集超过内存 ⇒ 换页到磁盘（thrashing）；或者用”并行算法跑在 1 个核上”作为基准（低估了基准）。CS149 版本明确警告：”把并行程序的加速比与’并行算法在单核上运行’相比，是常见且不诚实的做法。”

4.7 模型总览与适用条件

5. 关键要点

1. 先分清”延迟”与”带宽”，再决定优化手段：延迟问题（一次通信要等多久）靠减少消息数、增大消息体量（摊薄 T0）、重叠（overlap）、以及硬件改进；带宽问题（单位时间能搬多少）只能靠减少字节数（提高算术强度、消灭人为通信）。把这两件事搞混，优化就会南辕北辙。
2. 通信量由”工作分配”和”数据布局”共同决定，且必须先在算法层面压低下界：固有通信由分解/分配决定（1D blocked N/(2P)、1D interleaved 恒为 1/2、2D blocked N/(4√P)），人为通信由 cache line 粒度、write-allocate、放置策略、伪共享决定。好的分配（2D 分块）+ 好的布局（blocked layout + padding）能同时压两头。
3. 提高算术强度是并行程序最重要的单一优化目标：现代机器的算力/带宽之比很高（本讲示例机器 ridge point ≈ 38 flop/B），绝大多数”看起来很并行”的代码（逐元素运算、模板、稀疏计算）都落在带宽受限区，融合循环、分块重用、共享数据、合并消息都能把工作点往右推；只有当工作点越过 ridge point，计算优化才开始有意义。
4. 竞争是一种”与工作量无关”的性能税，要靠复制、细粒度化与错峰来缴清：吞吐上限 = 1/t_crit 与线程数无关，所以热点会让程序负扩展；有效手段是复制资源（本地副本、每桶锁、部分结果 + 合并）、把同步降级到片上/块内、用数据并行 + 排序取代细粒度同步，以及错开访问时间。
5. 优化的顺序必须是”先测量、先建 high watermark，再动手”（CS149 版本的方法论）：先写最简单、能跑对的并行版本并测出绝对时间；再用”加数学运算看是否线性变慢（指令受限？）”“把数组访问全改成 A[0]（局部性还有多少油水？）”“删掉所有原子/锁（同步还有多少油水？）”“删掉大部分计算但保留同样数据访问（是不是带宽受限？）”这四个实验确定瓶颈类型；只有在确认瓶颈类型之后，才去应用上面 4 条中的对应手段。

6. 常见陷阱与注意事项

• 把共享地址空间的”隐式通信”当成免费，同时只盯 FLOPS 不看字节：A[i] 这样一次普通 load 在并行程序里可能是跨核/跨插槽的一致性事务——”代码里没有 send/recv”不等于没有通信；反过来，把大量时间花在”减少浮点运算”上，而程序其实 90% 的时间在等内存，也是最常见的南辕北辙。先算算术强度并对照 ridge point，再决定优化计算还是优化访存。
• 伪共享（false sharing）：多个线程写同一条 cache line 的不同字节，逻辑上无共享、无数据竞争，硬件上却是”整条 line 的乒乓”（典型场景：count[tid]、struct { long a; long b; } 中不同线程各写一个字段、线程池的任务计数器）。修法：alignas(64) padding、per-thread 私有副本 + 末尾归约、blocked data layout。注意 padding 不是免费的（内存与带宽开销），对齐值要按机器的 cache line 大小（通常 64 B）定。
• 消息太小太频繁，以及非阻塞通信的四种误用：为每个元素发一条消息、每轮循环都做一次全局 Allreduce/屏障，都会被启动延迟 T0 碾压（应当合并消息、批量化、把归约放到低频路径上）；非阻塞通信的典型错误是 ① 在 Wait 之前就改写发送缓冲（数据被破坏）、② 忘记 Waitall、③ tag 不匹配（表现为永久挂起，看起来像死锁）、④ 用阻塞 send/recv 却让所有人都”先 send 再 recv”（真死锁，修法是奇偶配对收发或改用非阻塞）。
• 锁的粒度选错：全局锁 ⇒ 热点（吞吐被 1/t_crit 锁死，负扩展）；无脑细粒度锁 ⇒ 锁本身的开销（空间、原子操作、cache line 迁移）可能超过收益。先想清楚”临界区保护的到底是什么”，再决定粒度；能用”局部复制 + 末尾合并”或”数据并行 + 排序”就不用锁。
• 分块参数按错了 cache、以及忽略 NUMA 与数据放置：并行程序里每个线程只能用自己那部分 cache（以及自己那个 L3 slice），块大小必须按每线程可用容量调（太大装不下、太小则循环开销占比过高）；在多路机器上内存延迟与带宽还取决于”谁访问谁的内存”，而首次触碰（first touch）决定页的归属——应当在并行初始化时就让将来使用该数据的线程去初始化它。
• 用”固定小问题”测出的加速比评价程序或机器：小问题上通信/同步/启动开销占主导（通信/计算比高达 2P/N、T0 摊不薄），常常”不升反降”；同时要警惕用并行算法单核运行作基准造成的假加速，以及工作集在单机溢出 cache/内存（甚至换页到磁盘）造成的超线性加速——超线性往往说明基准跑得不对，而不是并行真的”多核快过 P 倍”。

7. 思考题（带答案）

思考题 1：三种工作分配的固有通信

设 N × N 网格、P = 64 个处理器、5 点模板（每个输出元素读上下左右四个邻居），N = 4096。请分别算出 1D blocked、1D interleaved（隔 P 行取一行）、2D blocked（8×8 处理器阵列） 三种分配下：每处理器计算的元素数、通信的元素数、以及”每通信 1 个元素对应多少计算元素”（算术强度），并说明哪种分配会随 N 增大而自动变好。

【答案】 设 P = 64、N = 4096，每处理器计算量都是 N²/P = 4096²/64 = 262,144 个元素。

• 1D blocked（连续 64 行）：通信量 = 上下各一行 = 2N = 8,192 个元素。 算术强度 = 262,144 / 8,192 = 32（公式 N/(2P) = 4096/128 = 32 ✓）。
• 1D interleaved（隔 64 行取一行）：每处理器拥有的行彼此间隔 64 行，每一条自己的行都要向邻居另要上下两整行（2N = 8,192 个元素），共有 N/P = 64 条这样的行 ⇒ 通信量 = 64 × 8,192 = 524,288。 算术强度 = 262,144 / 524,288 = 0.5（与 N、P 无关的常数 1/2 ✓）。比 1D blocked 差 64 倍，正好等于 P。
• 2D blocked（8×8 = 64 个处理器，每块 512×512）：通信量 = 四条边 = 4 × 512 = 2,048 个元素。 算术强度 = 262,144 / 2,048 = 128（公式 N/(4√P) = 4096/(4×8) = 128 ✓）。
• 随 N 变好的只有 blocked 两类：1D blocked 的算术强度 N/(2P) ∝ N；2D blocked 的 N/(4√P) ∝ N；而 1D interleaved 恒为 1/2，无论问题多大都注定是”搬 2 个元素才算 1 个元素”的灾难。此外 2D blocked 关于 P 是 ∝ 1/√P（次线性恶化），1D blocked 是 ∝ 1/P（线性恶化）——这就是讲义说”2D blocked 的通信扩展性渐近更好、捕获了算法的二维局部性”的定量含义。
• 补充（越大越好还是越小越好）：如果按讲义另一页的写法用通信-计算比（communication-to-computation ratio），三种分别是 1/32、2、1/128——越小越好，与上面的算术强度互为倒数，别被方向搞混。

思考题 2：伪共享的诊断与修复

某程序有 16 个线程，每个线程维护自己的局部统计量并只写自己的那个元素：

struct Stats { long hits; long misses; };
Stats stats[16];            // 16 * 16 B = 256 B，跨 4 条 cache line
...
#pragma omp parallel for
for (int i = 0; i < N; ++i) {
    int t = omp_get_thread_num();
    if (hit(i)) stats[t].hits++; else stats[t].misses++;
}

请回答：(a) 这段代码有没有数据竞争？(b) 为什么它的性能会随线程数增加而变差？(c) 给出至少两种修复方案，并说明各自的代价；(d) 你如何用实验确认瓶颈是伪共享而不是”真正的共享竞争”或者”内存带宽”？

【答案】 (a) 没有数据竞争：线程 t 只写 stats[t]，不同线程写得不同对象，且没有任何线程读别人写的值——从 C++ 内存模型看这是良定义、无竞争的程序。 (b) 原因在 cache line 粒度的一致性维护：Stats 是 16 字节，一条 64 B cache line 里放着 4 个线程的数据。线程 t 写自己的 hits 时必须取得整条 line 的独占所有权，于是其他 3 个线程的副本失效；它们下一次写又要把 line 抢回去。16 个线程的数据只占 4 条 line，这 4 条 line 在全核之间来回”乒乓”，每次写的代价从 ~0.5 ns（L1 命中）变成几十~几百 ns（line 所有权迁移）。这解释了为什么线程越多越慢：竞争这条 line 的人更多，而且线程很可能分布在不同的 CCX/插槽上（§3.4 实测：8 线程 13.6× 的差距、4 线程 7.9× 的差距——差距本身随线程数增大）。 (c) 修复方案：

1. padding / 对齐：struct alignas(64) Stats { long hits, misses; };（sizeof == 64）让每个线程独占一条 line。代价：内存与 cache 占用从 256 B 涨到 1 KB；如果线程数很大（比如 256），1 KB 还算便宜，但如果结构体本身是 100 B，padding 到 128 B 也会浪费。
2. 本地累加 + 末尾归约：线程把统计量累加到寄存器/栈上的局部变量，循环结束后只写一次 stats[t]。代价：需要一次额外的归约（reduction 或手工合并），并且改变了”随时可读全局统计”的语义（外部观察者只有在归约后才能看到完整值）。
3. blocked data layout：如果是数组而非结构体数组（例如 long hits[16]），可以考虑把”同一线程要用的所有数据”聚成连续的 block（struct alignas(64) PerThread { ... }），本质上与 1 相同。 (d) 确认瓶颈的实验（对应 CS149 版本的 high watermark 方法）：
   • 跑”删掉同步”实验：把写操作全部去掉（只做计算与读 ro[]），如果时间几乎不变 ⇒ 瓶颈不在同步，而是在数据搬运/cache line 迁移上 ⇒ 支持伪共享假设。§3.4 的版本 ③ 就是这个实验。
   • 跑”改数据布局”实验：只把 Stats 改成 alignas(64)，其他一字不改。如果加速 5~20 倍，伪共享被证实（真共享竞争不会因为 padding 而消失——padding 之后线程之间根本没有共享数据）。
   • 区分”带宽”：用 perf stat（或 Intel PCM / PAPI）看 LLC-load-misses、offcore/local vs remote DRAM 的字节数。伪共享的特征是DRAM 字节数很低、但一致性事务/延迟很高（数据本来就在 cache 里，只是反复易主）；真正的带宽瓶颈则是 DRAM 字节数接近机器带宽上限。
   • 注意：不要用”总时间随线程数恶化”作为唯一证据——带宽瓶颈与 NUMA 效应也会造成类似现象，必须结合字节数与 cache line 迁移计数。

思考题 3：三种分桶策略的选择

要把 100 万个粒子放进 256 个桶，机器是 16 线程的 CPU（量级参数：全局锁下一次插入 ≈ 600 ns，每桶锁下 ≈ 290 ns，数据并行分桶的吞吐 ≈ 每次插入 1.1 ns，但需要额外搬运 20n 字节；机器流式带宽 20 GB/s）。请估算三种方案的时间并给出选择；然后说明如果把目标机器换成 GPU（每 SM 可挂 2048 条线程、共 16 个 SM），你的选择会怎样改变，为什么。

【答案】

• 核算：n = 10⁶，额外字节 = 20n = 20 MB。
  • 全局锁：10⁶ × 600 ns = 0.60 s（且与线程数无关：加更多线程只会让队列更长，§3.5 实测 16 线程时吞吐只有 1.58 ~ 2.29 M/s，与此一致）。
  • 每桶锁：10⁶ × 290 ns = 0.29 s。虽然竞争降低，但插入路径上仍要原子操作 + 触碰桶头指针所在 line，没有解决”同步在插入路径上”这个根本问题。
  • 数据并行（计数排序 + 私有输出区间 + 并行求桶边界）：同步成本 ≈ 0（只有两次屏障），时间 ≈ 20 MB / 20 GB/s = 1 ms，加上直方图与散射本身的按元素开销（每次插入 1.1 ns 是实测含全部三个阶段的数字）⇒ 约 1~5 ms。比每桶锁快约两个数量级。
  • 选择：在这个规模上选数据并行（解法 5）。理由是：桶数（256）远小于线程数 × 常量、竞争系数极高，而额外带宽（20 MB）相对机器带宽很便宜。只有当桶数极少（比如 2 个）、线程数也极少（比如 2 个）时，额外的排序带宽才可能超过竞争成本，此时”每桶一把锁”或”局部复制 + 合并”更划算。
• 换成 GPU（16 个 SM，每个 SM 最多 2048 条线程 ⇒ 可并发数万条线程）：
  • 全局锁方案彻底不可行：热点吞吐上限是 1/t_crit，与线程数无关；数万条线程争一把锁只会把绝大部分时间花在排队上（而且 GPU 上原子操作/锁的成本比 CPU 更高，且 warp 内的锁会导致严重分歧与占用浪费）。
  • 每桶锁方案同样不可行：即使降到 256 把锁，每把锁平均也有上百条线程争抢；而且 GPU 的同步原语（atomicCAS 自旋）会占用 warp 调度槽位，属于”用宝贵的并发度换取等待”。
  • 应当选解法 ④ 或 ⑤：④（每个 thread block 维护一张局部网格，块内用 shared memory 同步，最后合并）把同步降级到片上（延迟比 L2/DRAM 低一个数量级，且不产生 cache line 迁移），代价是 N 倍内存与一次合并；⑤（数据并行：算格子号 → 排序 → 求桶边界）保留极大并行度且零细粒度同步，代价是一次排序与额外的数据遍历（额外带宽）——在 GPU 上 GPU 显存带宽很高（数百 GB/s 到 TB/s 级），额外带宽比”锁竞争”便宜得多，因此 ⑤ 往往是最优的；工程上还常把块内前缀和（__syncthreads() + shared memory）与块间偏移（全局前缀和）结合，做”块级直方图 + 块间合并”，这与解法 ④ 的思想是一致的。
  • 一般原则：机器的并发度越高，”用带宽/内存换取零同步”的方案就越划算；并发度越低、同步越廉价，”就地加锁”才越可能赢。 这正是讲义把五种解法并列出来的用意——没有普遍最优解，只有与机器匹配的解。

本讲与前后讲的关系回顾：本讲给出了”通信为什么贵、怎样变便宜”的完整工具箱（延迟/带宽 → α-β 模型 → overhead/occupancy/overlap → 固有 vs 人为通信 → blocking/fusion/sharing/granularity → 竞争与排队）；其中伪共享、cache line 迁移、一致性事务的硬件细节是第 11~13 讲（缓存一致性）的主题；网络拓扑如何决定 occupancy 与竞争是第 10 讲（互连网络）的主题；锁与原子操作的成本、无锁数据结构如何降低竞争是第 15~16 讲的主题。第 14 讲（Performance Analysis / Profiling）则会系统化本讲末尾提到的”测量与 high watermark”方法论。