Lecture 4: Parallel Programming Models

目录 · ← l3 · l5 →

Lecture 4: Parallel Programming Models

1. 章节标题与概述

Lecture 4: Parallel Programming Models

• 本讲核心问题：什么是”并行编程模型（programming model）”？它和”硬件实现（implementation）”有什么区别？讲义把这条主线概括为一句话——“Abstraction vs. implementation：把抽象与实现混为一谈，是这门课里最常见的困惑来源。” 本讲用三种通信抽象（共享地址空间 shared address space、消息传递 message passing、数据并行 data parallel）与三种机器架构的对照，回答”程序员该怎样思考并行程序的结构”以及”这些抽象背后需要什么硬件支持、代价是什么”。讲义正文标题页把它称作 “Lecture 3: Parallel Programming Abstractions (and their corresponding HW/SW implementations)”（见配套材料说明），前一半以 ISPC 为例讲 SPMD/SIMD，后一半讲三大通信模型。

• 涉及的主要硬件/软件机制：
  • 软件侧：ISPC 的 gang / programCount / programIndex / uniform / foreach、SPMD（单程序多数据）编程抽象、pthreads 线程、MPI 消息收发（send/recv + tag）、stream/gather/scatter 数据并行原语、以及编译器与并行运行时（compiler & parallel runtime）。
  • 硬件侧：SIMD 向量指令（SSE4 / AVX、AVX2 的 gather、AVX512 的 scatter）、SMP（”dance-hall” 组织）里的共享总线 / 多级网络 / 交叉开关（crossbar）、NUMA、片上网状/环状互连、fat tree 与 dragonfly 集群拓扑，以及”多核 + 多线程 + 超标量 + SIMD”的现代单芯片结构。
  • 接口侧：本讲反复强调的”系统接口”——OS 系统调用 API（如 pthread_create()）、ISPC 编译器产物（带 SIMD 指令的 .o）、MPI 库、以及数据并行语言（ISPC/OpenCL/CUDA）的 kernel 语义。

• 在并行计算知识体系中的角色：上一讲（A Modern Multi-Core Processor）讲的是”硬件给了你什么并行能力”（SIMD、多核、多线程、缓存层次）；本讲把它翻过来讲”软件用什么抽象去使用这些能力”。它是全课程的概念中枢：之后讲 CUDA（数据并行 + 同一 core 内共享地址空间）、OpenMP（共享地址空间 + fork-join）、MPI（消息传递）、缓存一致性（共享地址空间的硬件前提）、同步与内存序（共享地址空间的正确性代价），都要回到本讲的三个模型与”抽象距离（abstraction distance）”这个判据。讲义最后的结论也很明确：“实践中你必须能用多种方式思考”，因为现代机器在不同尺度上提供不同类型的通信，不同模型在不同尺度上最贴合机器。

• 配套材料：
  • lectures/04_progmodels.pdf（对应抽取文本 extracted/04_progmodels.txt，共 63 页幻灯片）：已公开，可在 https://www.cs.cmu.edu/~418/lectures/ 下直接下载（课程主页 https://www.cs.cmu.edu/~418/，日程表 https://www.cs.cmu.edu/~418/schedule.html，Fall 2026 日程中本讲位于 Aug 31，主题写作 “Parallel Programming Models”）。注意两个如实的细节：(a) 讲义第 1 页标题写作 “Lecture 3: Parallel Programming Abstractions”，与日程表中第 4 讲的主题名 “Parallel Programming Models” 略有差异，属讲义的编号/标题沿用现象；(b) 讲义页脚沿用历史学期（如 Fall 2025）字样，也属正常的讲义复用现象，不是错误。
  • cs149_supp/proghardware.txt：Stanford CS149 Fall 2025 Lecture 11 “Programming Specialized Hardware for AI”（共 60 页）——公开的补充读物。它与本讲的接口在于”编程模型如何反映硬件能力”这条主线：TPU 的 systolic array（脉动阵列）把”低控制开销 + 高数据复用”做成硬件，代价是编程模型变成”编译期就知道的矩阵乘/卷积 + 激活”这几条关键指令；SambaNova SN40L 的可重构数据流（reconfigurable dataflow）把编程模型变成 map/zip/reduce/gather/scatter 的数据流图 + metapipelining；NVIDIA H100/B100 则把异步（TMA、TMEM、异步 MMA）暴露给程序员，逼出了 ThunderKittens 这类嵌入式 DSL。注意：这部分属于补充视角，不是 CMU 15-418 讲义的原文内容，本笔记中出现处都会显式标注。
  • 讲课录像（Panopto / YouTube）：Fall 2026 日程表中被注释隐藏，属 未发布。
  • Ed 讨论区、Autolab、Canvas：需登录，非公开。
  • 历史学期 PDF（如 Performance Analysis/Profiling、Transactional Memory、AI in System Design 等 Fall 2026 尚未发布的讲义）位于 /afs/cs/academic/class/15418-*/public/ 之下，需要 CMU 登录，属未公开。
  • 本讲 Fall 2026 授课教师为 Brian Railing 与 Dimitrios Skarlatos；课程由 Kayvon Fatahalian 创建。

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

2.1 抽象 vs 实现：本讲的总纲

• 定义与目的：抽象（abstraction） 是”程序员看到并据以思考的东西”，实现（implementation） 是”硬件/运行时真正做的事情”。二者可以差距很大，而性能差异恰恰藏在差距里。讲义明确说：”conflating abstraction with implementation is a common cause for confusion in this course.”
• 直观解释（”它是什么？”）：把编程模型想成餐厅菜单，把实现想成厨房。菜单（抽象）写”宫保鸡丁”，你可能以为厨师是现炒的（顺序执行）；实际上厨房可能预先备好料、三位厨师并行分工、甚至在客人点单前就把热门菜下锅（乱序、推测、预取）。点菜的人（程序员）看到的是菜单，吃到的是厨房做出来的菜——菜单不保证速度，厨房决定速度。另一种类比：交通规则 vs 道路施工。规则（抽象）说”靠右行驶”；道路实际是几条车道、有没有立交（实现）才决定通行时间。
• 本讲的两个”抽象/实现”实例（讲义用它做全讲的开场）：
  1. ISPC：程序员看到的是 SPMD 抽象（gang 里有 programCount 条逻辑指令流），实现是一条 SIMD 指令流（AVX 的 8 个 lane）。
  2. pthreads：程序员看到的是”线程”这个抽象，实现是 OS 把线程映射到硬件执行上下文（execution context）（讲义最后一页的思考题：”who is responsible for mapping your pthreads to the processor’s thread execution contexts? Answer: the operating system”）。

2.2 系统分层：编程模型在哪里

• 定义与目的：讲义用一张分层图（”System layers: interface, implementation, interface, …“）说明每一层都是”上一层的实现、下一层的接口”。编程模型（programming model） 位于最上面几层：它是”描述并发/并行/独立计算的抽象”加上”描述通信的抽象”，为程序员提供思考程序结构的方式。
• 直观解释（”它是什么？”）：像洋葱/套娃：应用层说”我要并行”，语言/库把这句话翻译成原语，编译器与运行时把原语翻译成 OS 系统调用与机器指令，OS 把线程映射到硬件上下文，硬件再用流水线/SIMD/多核真正执行。每一层都只暴露一个”接口”给上层，把自己的”实现”藏起来。

图 1：系统分层（讲义 slide “System layers: interface, implementation, interface, …“）

   ┌──────────────────────────────────────────────────────────────────────┐
   │                        Parallel Applications                         │  ← 应用
   ├──────────────────────────────────────────────────────────────────────┤
   │  Language or library primitives / mechanisms      〔编程模型〕        │
   │    · 描述 concurrent / parallel / independent computation 的抽象      │
   │    · 描述 communication 的抽象                                       │
   ├──────────────────────────────────────────────────────────────────────┤
   │                Compiler and/or parallel runtime                      │
   ├──────────────────────────────────────────────────────────────────────┤
   │                       Operating system                               │
   │            OS system call API  ← Red italic: 系统接口                │
   ╞══════════════════════════════════════════════════════════════════════╡
   │              Hardware Architecture  ← HW/SW boundary                 │
   ├──────────────────────────────────────────────────────────────────────┤
   │         Micro-architecture (hardware implementation)                 │
   └──────────────────────────────────────────────────────────────────────┘
     Blue italic = abstraction/concept   Red italic = system interface
     Black       = system implementation
   关键操作/性能特征：
     · 每跨一层都要"翻译"：抽象语义不会被实现原样保留 ⇒ 性能不可从程序文本直接读出
     · "abstraction distance"（抽象距离）越小 ⇒ 性能越可预测，但代码可移植性/灵活性越差

以 pthreads 为例，这张图会具体化为：Parallel Application → 抽象”thread” → pthread 库实现 → pthread_create()（系统调用 API）→ OS 内核线程管理 → x86-64 现代多核 CPU。以 ISPC gang 为例则变成：Parallel Application → ISPC 语言（调用 ISPC 函数、foreach 构造）→ ISPC 编译器 → x86-64（含 AVX 向量指令）→ CPU 的单核。注意后者：讲义特别提示，”仅 gang 抽象而言，这一切只跑在一个核上”；多核执行要靠 ISPC 的另一个原语 task。

2.3 ISPC 与 SPMD：gang、programCount、programIndex、uniform

• 定义与目的：ISPC（Intel SPMD Program Compiler） 实现 SPMD（single program multiple data，单程序多数据） 抽象。调用一个 ISPC 函数会生成（spawn）一个 gang，gang 中的每条逻辑指令流称为一个 program instance；所有 instance 并发运行同一段 ISPC 代码，返回时所有 instance 一起完成。
  • programCount：gang 中同时执行的 instance 数（一个 uniform 值，所有 instance 相同）。
  • programIndex：当前 instance 在 gang 中的编号（一个 varying 值，各 instance 不同）。
  • uniform：类型修饰符，表示”所有 instance 值相同”。讲义原话：“Its use is purely an optimization. Not needed for correctness.”（纯粹是优化手段，正确性不依赖它）——但它是让实现细节”从抽象里透出来”的关键，例如 uniform int denom 可以让分母在向量化后只存一份。
  • foreach (i = 0 ... N)：声明”这些迭代由 gang 中的 instance 协作完成”，即 map（映射）；ISPC 当前实现采用静态交错（interleaved）分配，但抽象本身允许其他分配方式。
• 直观解释（”它是什么？”）：把 gang 想成合唱团（chorus）：programCount 是团员人数，programIndex 是团员编号，他们共用同一份乐谱（同一指令流）、但在各自的声部/歌词（各自的数据）上唱。指挥（编译器）把乐谱翻译成”一条 SIMD 指令”，八个团员在同一拍上各唱自己的词。而 uniform 就像”全体一起唱的那句”——大家词一样，可以只看一份谱。另一种类比：同一个操练口令下的队列。
• 性能要点（抽象 vs 实现）：程序员”认为”自己 spawn 了 programCount 条指令流；编译器把 gang 实现为 SSE4/AVX 向量指令，并用掩码（masking） 处理条件控制流。gang 的大小 = 硬件的 SIMD 宽度（或 SIMD 宽度的小整数倍）。这带来两个直接结论：(1) 同一个 gang 里的 instance 是同步推进的，因此它们之间不存在”交错执行”；(2) 条件分支若不同 lane 走不同路径，就会串行化（谓词化/掩码）——这是后面性能分析的关键。

图 2：ISPC 的 SPMD 抽象与 SIMD 实现（软件执行模型）

 抽象层（程序员看到的世界）           实现层（ISPC 编译器生成的东西）
 ─────────────────────────           ────────────────────────────────────────
  main.cpp: 顺序 C++ 代码
      │
      │  sinx(N, terms, x, result)          一条 SIMD 指令流（AVX, 8 lane）
      ▼                                     ┌──────────────────────────────┐
 ┌──────────────────────────────┐           │ vload  x[i .. i+7]           │
 │  gang: programCount = 8      │           │ vmul   x*x*x                 │
 │  ┌────┬────┬────┬────┐       │  ═══►     │ vadd/vdiv (terms 次)         │
 │  │ i0 │ i1 │ i2 │ i3 │       │  编译     │ vstore result[i .. i+7]      │
 │  ├────┼────┼────┼────┤       │           └──────────────────────────────┘
 │  │ i4 │ i5 │ i6 │ i7 │       │            lane0=i0 … lane7=i7
 │  └────┴────┴────┴────┘       │            8 个 lane 在同一拍推进（锁步）
 │  8 条逻辑指令流（不同 programIndex）       条件控制流 → 掩码（mask）实现
 └──────────────┬───────────────┘
                │ 所有 instance 完成后一起返回（gang 返回 = 隐式屏障）
                ▼
            恢复顺序执行（C++）
   性能特征：
     · 一条 packed load/store 覆盖 8 个 float = 32 B，一次访问即满足整个 gang
     · 吞吐 ≈ 每时钟一条 8 宽 SIMD 指令；延迟被多线程/乱序隐藏（见上讲的处理器结构）
     · 若 8 个 lane 的访存不连续 → 退化为 gather（代价高，见 2.4）

2.4 交错分配 vs 分块分配：packed load vs gather

• 定义与目的：在 SPMD 里”哪个 instance 处理哪些迭代”直接决定访存模式，进而决定用哪条 SIMD 指令。
  • 交错分配（interleaved）：第 t 轮里，instance k 处理元素 idx = t*programCount + k。相邻 lane 访问相邻地址→ 一条 packed load（讲义举 _mm_load_ps1）即可完成全部 instance 的取值。这是 foreach 默认的行为。
  • 分块分配（blocked）：instance k 处理 [k*count, (k+1)*count)。第 t 轮里各 lane 访问地址相差 count（跨步 stride）→ 必须用 gather（讲义举 _mm_i32gather），是”更复杂、更昂贵”的指令。
• 直观解释（”它是什么？”）：交错像发牌——一张张轮流发给四个人，每轮四个人手里的牌凑起来正好是连续的一段（一条 cache line 就够）；分块像切蛋糕——每个人抱走一大块，四个人同时伸手去拿的是相隔很远的四块，得跑四趟。发牌（交错）对缓存与向量单元都友好；切蛋糕（分块）适合”每人一份独立任务”的直觉，但对 SIMD 不友好。
• 讲义给出的硬件时间线（很有用的事实）：gather 指令 2013 年随 AVX2 才有；AVX2 不支持 SIMD scatter（scatter 只能用标量循环模拟）；scatter 指令出现在 AVX512；GPU 上硬件支持的 gather/scatter 是存在的，但相对连续向量的 load/store 仍然昂贵。

图 3：交错 vs 分块的访存模式（programCount = 4，处理 16 个元素）

  交错 interleaved（foreach 默认）               分块 blocked（手写索引）
  idx = i + programIndex                        idx = start + i, start = programIndex*count
  ┌────────────────────────────────┐            ┌────────────────────────────────┐
  │ t=0: lane0..3 → 元素 0,1,2,3   │ 连续 16B   │ t=0: lane0..3 → 元素 0,4,8,12  │ 跨步 4*4B
  │ t=1: lane0..3 → 元素 4,5,6,7   │ 连续 16B   │ t=1: lane0..3 → 元素 1,5,9,13  │ 跨步 4*4B
  │ t=2: lane0..3 → 元素 8,9,10,11 │ 连续 16B   │ t=2: lane0..3 → 元素 2,6,10,14 │ 跨步 4*4B
  │ t=3: lane0..3 → 元素 12..15    │ 连续 16B   │ t=3: lane0..3 → 元素 3,7,11,15 │ 跨步 4*4B
  └────────────────────────────────┘            └────────────────────────────────┘
       一条 packed load / store                       需要 gather / scatter
       （_mm_load_ps1 / 向量 store）                 （_mm_i32gather；AVX2 无 scatter）
   每条 packed 访存搬 4×4B=16B，1 次事务           每 lane 独立地址 ⇒ 多条事务、更高延迟

表：两种静态分配的对比（本讲核心工程结论之一）

2.5 通信抽象之一：共享地址空间（shared address space）

• 定义与目的：线程通过读写共享变量通信；通信是隐式的（隐含在 load/store 里）；同步原语（lock、semaphore）本身也是共享变量。讲义定位它是顺序编程的自然延伸：”In fact, all our discussions in class have assumed a shared address space so far!”
• 直观解释（”它是什么？”）：讲义给的类比是一块大公告板（bulletin board）：谁都可以往上贴，谁都可以看；线程 1 写 X，稍后线程 2 读 X 就”看到”了更新。更生活化的补充类比：合租公寓的公共冰箱 + 便利贴——方便，但没人规定”谁先拿”、”你看到的是不是最新那张”，需要额外的规则（锁）才能保证不出乱子。这就是共享地址空间的根本特征：极其自然，但也极其容易”搬起石头砸自己的脚”（讲义原话：natural way of programming, but can shoot yourself in the foot easily）。
• 硬件实现（关键点）：任意处理器都能直接引用任意内存地址。典型组织是 SMP（symmetric multiprocessor）”dance-hall” 结构：处理器各自带本地 cache，通过互连（总线、多级网络、交叉开关）访问集中式内存。讲义强调一个前提：”cost of accessing an uncached memory address is the same for all processors”（caching 会引入非均匀访问时间，这正是后面缓存一致性讲座的内容）。性能特征：SMP 的优点是”成本统一”，缺点是”统一地差”（memory is uniformly far away）——所以 SMP 不可扩展，规模一大就必须转向 NUMA 或集群。

图 4：共享地址空间的硬件实现（SMP “dance-hall” 与 NUMA）

 (a) SMP / "dance-hall"：所有处理器等距访问同一份内存
   ┌──────────┐ ┌──────────┐ ┌──────────┐ ┌──────────┐
   │Processor │ │Processor │ │Processor │ │Processor │
   │ L1/L2 $  │ │ L1/L2 $  │ │ L1/L2 $  │ │ L1/L2 $  │
   └────┬─────┘ └────┬─────┘ └────┬─────┘ └────┬─────┘
        └────────────┴─────┬──────┴────────────┘
                     ┌─────┴──────┐   互连实现三选一：
                     │Interconnect│   · 共享总线 shared bus（便宜，带宽争用）
                     └─┬────────┬─┘   · 多级网络 multi-stage network
                       │        │     · 交叉开关 crossbar（无阻塞，面积/成本高）
              ┌────────┴──┐  ┌──┴───────┐
              │  Memory   │  │   I/O    │  任何处理器可直接引用任何地址
              └───────────┘  └──────────┘

 (b) NUMA（non-uniform memory access）：都能访问，但"远近不同"
   ┌───────────┐ ┌───────────┐          ┌───────────┐ ┌───────────┐
   │  CPU 1-4  │ │  Memory   │          │  CPU 5-8  │ │  Memory   │
   │  + L1/L2$ │ │ (本地/近) │          │  + L1/L2$ │ │ (本地/近) │
   └─────┬─────┘ └─────┬─────┘          └─────┬─────┘ └─────┬─────┘
         └──────┬──────┘                      └──────┬──────┘
                │                                    │
                └─────────────────┬──────────────────┘
                         ┌────────┴────────┐
                         │  片间/节点互连  │   AMD HyperTransport / Intel QPI·UPI
                         └─────────────────┘
   性能特征（讲义原话的意思）：
     · 保持"统一访问时间"的代价是 scalability 崩塌：GOOD = 成本统一，BAD = 统一地差
     · NUMA 更可扩展：本地内存低延迟、高带宽；代价 = 程序员必须"找局部性"
     · 例：访问地址 X，从 core 5-8 的延迟显著高于 core 1-4

讲义列出的真实机器（都属”共享地址空间”这一类家族）：Intel Core i7 四核（片内互连是 ring）、AMD Phenom II 六核、SUN Niagara 2 (UltraSPARC T2) 八核（crossbar 交叉开关，面积约等于一个核）、SGI Altix UV 1000（256 blade × 2 CPU × 8 核 = 4096 核，单一共享地址空间，互连为 fat tree）、以及 LLNL El Capitan（11,424/11,520 节点异构集群，节点内 4× AMD Instinct MI300A APU，每 APU 24 核 Zen 4 + 228 CU，统一地址空间 CPU+GPU，5.3 TB/s，节点间 Infinity Fabric 缓存一致 all-to-all + Slingshot 200 Gb/s NIC，网络拓扑为 dragonfly）。

2.6 通信抽象之二：消息传递（message passing）

• 定义与目的：线程各自拥有私有地址空间，只能靠收发消息交换数据：send 指定接收者、要发的缓冲区、可选的消息标识（tag）；recv 指定发送者、存放数据的缓冲区、可选的 tag。发送消息是两个线程之间交换数据的唯一途径。最流行的软件库是 MPI（message passing interface）。
• 直观解释（”它是什么？”）：像两个不通气的工作间之间的传递窗/快递：你没法直接伸手去拿别人的东西，必须打包、贴单号（tag）、投递，对方凭单号签收。好处是账目清楚——读代码就能看出”什么时候发生了通信”（讲义：can read program and see where the communication is）；坏处是第一次写对更难（要自己管缓冲、匹配、分区），但结构上的强制约束往往帮你更快得到可扩展的正确程序。
• 硬件实现与性能特征：硬件不必实现全系统范围的 load/store，只要能通信即可——所以可以用商品机 + 网络（讲义例子：Infiniband 上的 workstation 集群、IBM Blue Gene/P 超算）。关键性能事实：当消息传递跑在”硬件其实是共享地址空间”的机器上时，“发消息”就是内存拷贝（从消息库缓冲区拷出/拷入），这决定了它的开销模型是”延迟 + 带宽”（详见第 4 节）。

图 5：消息传递模型（两个私有地址空间 + 唯一通道）

   Thread 1 的地址空间                      Thread 2 的地址空间
  ┌───────────────────────────┐          ┌───────────────────────────┐
  │  变量 X  ┌──────────┐     │          │  变量 Y  ┌──────────┐     │
  │          │  value   │     │          │          │   ???    │     │
  │          └────┬─────┘     │          │          └────┬─────┘     │
  │               │ send(X,2,tag)        │     recv(Y,1,tag) │       │
  └───────────────┼───────────┘          └─────────────────┼─────────┘
                  │                                        │
                  ▼          消息（库缓冲区/网络）           ▲
        ┌──────────────────────────────────────────────────────────┐
        │  send: 接收者 + 缓冲区 + tag     recv: 发送者 + 缓冲区+tag│
        │  ⇒ 数据被"拷贝"过网络，"通信点"在源代码中显式可见         │
        └──────────────────────────────────────────────────────────┘
   性能特征：
     · 只有两种成本来源：每条消息的固定延迟 α（启动/协议/拷贝） + 每字节成本 1/β
     · 无 cache line 级别的隐式通信 ⇒ 无伪共享问题；但每次通信都要显式拷贝
     · 可扩展到"共享地址空间在硬件上无法维持"的规模（集群/超算）

2.7 通信抽象之三：数据并行（data parallel）与 stream 模型

• 定义与目的：对集合（collection）中的每个元素施加同一个函数，且不同元素之间相互独立。历史形态：80 年代 SIMD 巨型机（Connection Machine CM-1/CM-2：数千处理器 + 单一指令译码单元）、Cray 向量处理器（add(A,B,n) 是”一条指令处理长度为 n 的整个向量”）、Matlab（C = A + B）。今天的常见形态是 SPMD + map：map(function, collection)，函数作用在每个元素上，同步隐含在 map 结束处（map 返回时函数已作用到所有元素）。
• 直观解释（”它是什么？”）：像流水线盖章：一叠纸每张都盖同一个章，纸与纸之间不需要商量；也像考场上的同一份卷子：每个考生（元素）做同一套题（同一个函数），彼此不通信（”no communication among distinct function invocations”）。map 的语义允许实现以任意顺序、包括并行地调度各次调用——这就是数据并行”能自动并行”的根本原因。
• ISPC 里的数据并行：foreach 就是 map，循环体就是 function。但讲义提醒一个关键细节：“collection” 在 ISPC 里不是一等概念——它由程序里的数组索引逻辑隐式定义，ISPC 里没有“把这个代码 map 到这个数组的所有元素上”这种语义的操作。三个例子清楚展示了这一点：
  • absolute_value：y[i] = \|x[i]\| → 最标准的逐元素 map。
  • absolute_repeat：y[2i] = y[2i+1] = \|x[i]\| → 仍然是合法 ISPC 程序，但”把循环体 map 到已有集合上”的直觉就变得别扭（一个迭代写两个输出位置）。
  • shift_negative：if (i>=1 && x[i]<0) y[i-1] = x[i]; else y[i] = x[i]; → 这个程序是非确定的（non-deterministic）！ 多个迭代可能写同一内存位置，而数据并行模型不规定迭代执行顺序，也不提供细粒度互斥/同步原语（它本就不是为这种结构设计的）。
• “更正统”的数据并行：stream 编程模型：显式定义集合（stream） 与无副作用函数（kernel），把每次调用的输入/输出/临时量视为私有地址空间。讲义（用自造的语法）给出：

  stream<float> x(N), y(N);   // 定义集合
  absolute_value(x, y);       // 把 kernel map 到集合上
  void absolute_value(float x, float y) { if (x<0) y = -x; else y = x; }
  

  收益（讲义明确列出）：函数真的无副作用（写不出非确定程序）；数据流对编译器已知（每次调用的输入输出提前可知 → 可以做 prefetch 预取隐藏延迟）；生产者-消费者局部性提前可知（可把前一个 kernel 的输出直接喂给后一个 kernel，中间值留在片内 buffer/cache、根本不写回内存，省带宽）。 代价（”正统”数据并行的阿喀琉斯之踵）：需要一套算子库来描述复杂数据流，而一旦数据流稍复杂，就得”cross fingers and hope compiler is intelligent enough”——讲义原话：”If I just had one more operator…“。

• 两个关键通信原语：gather / scatter。stream_gather(input, indices, tmp) 按索引收集，stream_scatter(tmp, indices, output) 按索引写散。ISPC 等价写法就是把 input[indices[i]] 或 output[indices[i]] 写进 foreach 里。

图 6：stream/数据并行的融合（kernel fusion）与片内传递

  朴素三步（每个 kernel 结束后中间结果落回 DRAM）：
     input ──►┌───────┐──► tmp（写 DRAM）──►┌───────┐──► output
              │  foo  │                      │  bar  │
              └───────┘                      └───────┘
     访存量 ≈ 2N(读 input) + 2N(写 tmp) + 2N(读 tmp) + 2N(写 output)

  融合后（编译器知道数据流 ⇒ 生产者-消费者局部性已知）：
     input ──►┌───────┐──►〔on-chip buffer / cache〕──►┌───────┐──► output
              │  foo  │        tmp 从不回 DRAM          │  bar  │
              └───────┘                                 └───────┘
     访存量 ≈ 2N(读 input) + 2N(写 output) ⇒ 省掉 4N 字节的 DRAM 往返

  等价的手写形式（讲义给出）：
     for (i=0; i<N; i++) output[i] = bar(foo(input[i]));
   性能特征：省的是【DRAM 带宽】而不是【算力】⇒ 对算术强度低的 kernel 效果巨大

2.8 模型与机器：对应关系是”模糊的（fuzzy）”

• 定义与目的：讲义用一整页强调：编程模型（描述程序的抽象）与机器类型（硬件的实现）之间不是一一对应。
  • 在”硬件实现共享地址空间”的机器上实现消息传递很常见：发消息 = 从消息库缓冲区拷贝内存；收消息 = 从库缓冲区拷回来。
  • 在”硬件不支持共享地址空间”的机器上，也能用（效率较低的）软件实现共享地址空间抽象：把含共享变量的页标记为 invalid，缺页异常处理程序（page-fault handler） 发出相应的网络请求把页取来（这属于共享虚拟内存/软件 DSM 的做法）。
  • 因此必须始终区分：”什么是编程模型（用来描述程序的抽象）？什么是硬件实现？”
• 直观解释（”它是什么？”）：像语言与口音。你可以用普通话（抽象）在电话里说（一种实现），也可以面对面说（另一种实现）；反过来，你也可以用英语（另一种抽象）在同一个电话里说。“说什么”和”怎么传”是两件独立的事。
• 工程推论（讲义总结页的原则）：三种模型施加的限制（restrictions） 是有意设计来”反映并行化与通信成本的现实”的：
  • 共享地址空间机器：硬件支持任意处理器访问任意地址；
  • 消息传递机器：硬件可能提供加速 send/recv/缓冲的机制；
  • “abstraction distance” 应当尽量小以让性能可预测，但要足够大以保证代码的灵活性与可移植性。

表：三种通信抽象的对比（本讲的核心表格）

2.9 单芯片尺度：程序如何落到”四核 + SMT + SIMD”上

• 定义与目的：讲义最后 5 页把前面的抽象收回到硬件：程序里的并行究竟落在哪个硬件资源上？它给出了一台”教学用处理器”的完整结构：4 个核，每核 2 个执行上下文（2-way multithreading），每时钟最多 2 条指令，其中一条是 8 宽 SIMD；每核有私有 L1/L2，共享 L3 与内存控制器。
• 直观解释（”它是什么？”）：把核想成厨房里的灶台，执行上下文（execution context）是同时备着的多口锅：一口锅等水开（内存延迟）时，厨师去照看另一口锅（硬件多线程隐藏延迟）；”每时钟 2 条指令”是两只手同时干活（超标量）；”8 宽 SIMD”是一次切八份同样的菜（数据并行）。而 OS 是排班经理，把程序员交来的线程（pthreads）分配给这些”锅/灶台”。

图 7：四核 × 2 执行上下文 × 8 宽 SIMD 的单芯片结构（讲义最后一页的结构）

 ┌─────────────────────────────── Chip ───────────────────────────────────┐
 │ ┌──────────────┐ ┌──────────────┐ ┌──────────────┐ ┌──────────────┐    │
 │ │  Core 0      │ │  Core 1      │ │  Core 2      │ │  Core 3      │    │
 │ │ ┌──────────┐ │ │ ┌──────────┐ │ │ ┌──────────┐ │ │ ┌──────────┐ │    │
 │ │ │ EC0 │ EC1 │ │ │ │ EC0 │ EC1 │ │ │ │ EC0 │ EC1 │ │ │ │ EC0 │ EC1 │ │    │
 │ │ └──┬───┬───┘ │ │ └──┬───┬───┘ │ │ └──┬───┬───┘ │ │ └──┬───┬───┘ │    │
 │ │  F/D  F/D    │ │  F/D  F/D    │ │  F/D  F/D    │ │  F/D  F/D    │    │
 │ │  [SIMD Exec2]│ │              │ │              │ │              │    │
 │ │  [Exec 1   ] │ │   ...同构... │ │              │ │              │    │
 │ │ ┌──────────┐ │ │              │ │              │ │              │    │
 │ │ │ L1 Cache │ │ │              │ │              │ │              │    │
 │ │ ├──────────┤ │ │              │ │              │ │              │    │
 │ │ │ L2 Cache │ │ │              │ │              │ │              │    │
 │ │ └────┬─────┘ │ │              │ │              │ │              │    │
 │ └──────┼───────┘ └──────────────┘ └──────────────┘ └──────────────┘    │
 │        └──────────────┬──────────────────────────────────────┘         │
 │                  ┌────┴─────┐   片内互连（如 ring）                     │
 │                  │ L3 Cache │                                          │
 │                  └────┬─────┘                                          │
 │              ┌────────┴────────┐                                       │
 │              │ Memory Contr.   │──── Memory Bus (to DRAM)              │
 │              └─────────────────┘                                       │
 └─────────────────────────────────────────────────────────────────────────┘
   规格：4 核 × 2 执行上下文（最多 8 个活跃线程）× 每核每时钟 ≤2 条指令
         （其中一条是 8 宽 SIMD 指令）
   性能特征：
     · 单核吞吐上限 = 1 条 8 宽 SIMD/时钟（数据并行）+ 1 条标量/时钟（指令级并行）
     · 内存延迟靠"另一个执行上下文"与乱序执行隐藏（延迟隐藏，不是带宽提升）
     · 4 核共享 L3 与 DRAM 带宽 ⇒ 多核扩展的上限常常是【带宽】而非【算力】
   讲义思考题（本讲结论的收口）：
     · 两个 pthread 如何映射到 4 个执行上下文？→ 由 OS 决定
     · 若 spawn 五个 pthread 呢？→ 出现"线程数 > 上下文数"的过订阅（oversubscription）

2.10 补充视角（Stanford CS149，非 418 讲义原文）：把编程模型”钉”在硬件能力上

• 定义与目的：CS149 的补充材料讲专用硬件（specialized hardware）如何反向塑造编程模型，与本讲”抽象必须反映硬件通信成本”的主题完全同构：
  • Google TPU：用 systolic array（脉动阵列） 高效做稠密矩阵乘。芯片面积中算术单元约占 30%，控制逻辑面积占比很低；关键指令只有寥寥几条：read host memory、write host memory、read weights、matrix_multiply/convolve、activate。也就是说，“编程模型”被压缩成了”喂数据 + 矩阵乘 + 激活”——极致效率的代价是极小的可编程域。
  • 脉动阵列 vs SIMD（讲义补充材料中的对比表）：SIMD 是控制驱动（instructions）、复用有限、通信走全局（寄存器/内存）、控制集中；systolic array 是数据驱动（wavefront 波前）、具有时间与空间复用、通信只在相邻 PE 之间（局部）、控制分布式。
  • 数据流（dataflow）编程模型：SambaNova SN40L 的 RDU 用 PCU（Pattern Compute Unit）+ PMU（Pattern Memory Unit）+ mesh 开关 组成可重构数据流阵列；编程模型是 map / zip / reduce / gather / scatter 等可组合原语，再经过 tiling → parallelization → metapipelining → place & route → codegen 落到硬件。Metapipelining（元流水线） = “pipeline of pipelines”：把循环体拆成流水级、级间用双缓冲（double buffer） 传递中间数据，从而重叠多个循环迭代并容忍各级耗时不均。
  • 异步（asynchrony）与 DSL：H100/B100 把 Tensor Core、TMA（Tensor Memory Accelerator）、TMEM、异步 MMA 暴露出来，编程复杂度飙升，于是出现 ThunderKittens / Mosaic / CuTe-DSL 这类嵌入式 DSL：设计原则是”16×16 tile 作为原始数据类型”、”处处异步”、”提供高层 GPU 协作模式（生产者-消费者）”。
• 直观解释（”它是什么？”）：这是“为特定活法定制的流水线车间”。通用 CPU 像万能工具箱（什么都能干，效率一般）；systolic array 像专门压铸某零件的自动线（快、省电，但只能干这一件事）；可重构数据流像可重排的乐高传送带，靠编译器当”工艺工程师”把计算摊在空间上。补充材料里的 “Hardware Lottery”（Sara Hooker）说得更直白：某个研究想法胜出，往往是因为它适配了当时的软硬件，而不是因为它普适地更优。

图 8（补充）：SIMD 执行 vs 脉动阵列（systolic array）数据流

 (a) SIMD（控制驱动，本讲前半部分的主角）
     ┌──────────── 取指/译码（集中）────────────┐
     │  vload x[i..i+7] / vmul / vstore         │  ← 同一指令流驱动全部 lane
     └───────┬────────┬────────┬────────┬───────┘
           lane0    lane1    lane2   ...  lane7     lane 之间不直接通信，
             │        │        │              │      数据经寄存器/内存"全局"交换
           [ALU]    [ALU]    [ALU]          [ALU]
     复用：有限（靠寄存器重用/缓存）；控制：集中；每 mm² / 每瓦效率：低

 (b) Systolic array（数据驱动，权重驻留 + 波前推进）
      x0 ──► ┌────┐ ──► ┌────┐ ──► ┌────┐ ──► ┌────┐
              │PE00│     │PE01│     │PE02│     │PE03│ ──► y0 (累加器)
      x1 ──► ┌────┐ ──► ┌────┐ ──► ┌────┐ ──► ┌────┐
              │PE10│     │PE11│     │PE12│     │PE13│ ──► y1
              └────┘     └────┘     └────┘     └────┘
       权重 FIFO：w00 w01 w02 w03 / w10 ... 沿行"注入"，驻留在 PE 内
     每个 PE：f = w*x + (来自上方/左方的部分和)，一拍一乘一加
     复用：时间 + 空间双重；控制：分布式；效率：高（面积几乎全给算术单元）
     代价：可编程域极窄（关键指令只有矩阵乘/卷积/激活 + 数据搬运）

3. 代码示例与性能分析

3.1 示例 1：ISPC 的 sin(x)——SPMD/SIMD 与交错/分块访存

讲义用同一个 sinx 函数贯穿前半讲：先给 C 版本，再给 ISPC（交错版）、ISPC（分块版）、以及用 foreach 提升抽象层次的版本。

// ============ 文件: sinx.ispc ============
// 编译（生成 SIMD 目标文件与 C 头文件）：
//   ispc --target=avx2-i32x8 -O3 sinx.ispc -o sinx_ispc.o -h sinx_ispc.h
// （-O3 为 release 优化；--target=avx2-i32x8 表示 AVX2、8 宽 32-bit lane，
//   即 gang 大小 programCount = 8）

export void sinx(uniform int N, uniform int terms,
                 uniform float x[], uniform float result[])
{
    // 交错分配（ISPC 的默认/foreach 语义）：交错访问相邻元素
    // 假设 N % programCount == 0
    foreach (i = 0 ... N) {
        float value = x[i];
        float numer = x[i] * x[i] * x[i];
        uniform int denom = 6;      // 3!   —— uniform：8 个 instance 共用一份
        uniform int sign  = -1;     // uniform：纯优化，不影响正确性
        for (uniform int j = 1; j <= terms; j++) {
            value  += sign * numer / denom;
            numer  *= x[i] * x[i];
            denom  *= (2 * j + 2) * (2 * j + 3);
            sign   *= -1;
        }
        result[i] = value;          // 一条 packed store 覆盖 8 个 float
    }
}

// ---- 等价的"手写交错版"（不用 foreach，自己算索引，讲义的做法）----
export void sinx_interleaved(uniform int N, uniform int terms,
                             uniform float x[], uniform float result[])
{
    for (uniform int i = 0; i < N; i += programCount) {
        int idx = i + programIndex;          // programIndex: 非 uniform（varying）
        float value = x[idx];
        float numer = x[idx] * x[idx] * x[idx];
        uniform int denom = 6;
        uniform int sign  = -1;
        for (uniform int j = 1; j <= terms; j++) {
            value += sign * numer / denom;
            numer *= x[idx] * x[idx];
            denom *= (2 * j + 2) * (2 * j + 3);
            sign  *= -1;
        }
        result[idx] = value;
    }
}

// ---- 分块版：每 instance 负责连续一段，访存变成跨步 ⇒ 需要 gather ----
export void sinx_blocked(uniform int N, uniform int terms,
                         uniform float x[], uniform float result[])
{
    uniform int count = N / programCount;    // 假设整除
    int start = programIndex * count;        // varying
    for (uniform int i = 0; i < count; i++) {
        int idx = start + i;                 // 各 lane 的 idx 相差 count ⇒ gather
        float value = x[idx];
        float numer = x[idx] * x[idx] * x[idx];
        uniform int denom = 6;
        uniform int sign  = -1;
        for (uniform int j = 1; j <= terms; j++) {
            value += sign * numer / denom;
            numer *= x[idx] * x[idx];
            denom *= (2 * j + 2) * (2 * j + 3);
            sign  *= -1;
        }
        result[idx] = value;
    }
}

// ============ 文件: main.cpp（顺序 C++ 侧）============
// 链接：g++ -O3 -mavx2 main.cpp sinx_ispc.o -o sinx
#include "sinx_ispc.h"      // ISPC 编译器生成
#include <cstdio>
#include <cstdlib>
#include <chrono>

int main() {
    const int N     = 1 << 24;      // 16M 个 float = 64 MB 输入
    const int terms = 5;

    float* x      = new float[N];
    float* result = new float[N];
    for (int i = 0; i < N; i++)                  // 初始化（略：填 [-1,1] 的随机值）
        x[i] = (float)((i % 1000) - 500) / 500.0f;

    auto t0 = std::chrono::steady_clock::now();
    sinx(N, terms, x, result);                   // 调用 ISPC 函数 ⇒ spawn 一个 gang
    auto t1 = std::chrono::steady_clock::now();

    double sec = std::chrono::duration<double>(t1 - t0).count();
    printf("N=%d terms=%d  time=%.3f ms\n", N, terms, sec * 1e3);
    printf("result[0]=%f  result[N-1]=%f\n", result[0], result[N - 1]);
    delete[] x; delete[] result;
    return 0;
}

【代码做什么？】

1. main.cpp 顺序执行：分配两个长度为 N 的 float 数组（共 128 MB），初始化 x。
2. 调用 sinx(N, terms, x, result)——这一步会 spawn 一个 gang（本例 8 个 program instance），8 条逻辑指令流并发进入 ISPC 函数体。
3. 函数体内 foreach (i = 0 ... N) 声明”这 N 个迭代由 gang 协作完成”；实现把迭代交错分给 instance：第 t 轮里 instance k 处理 idx = t*8 + k。
4. 每个迭代用泰勒展开算 sin(x[i]) ≈ x - x³/3! + x⁵/5! - x⁷/7! + x⁹/9!（terms=5）。
5. 8 个 instance 全部完成后 gang 一起返回（隐式屏障），main 恢复顺序执行。
6. 分块版 sinx_blocked 做同样的事，但把连续一段分给同一 instance：语义相同，访存模式完全不同。

说明（如实起见）：讲义里用 foreach 的那个版本，循环体内最后一行写作 result[idx] = value;，而 idx 是”手写交错版”才定义的变量；按 foreach 的语义此处应为 result[i] = value;。上面代码按语义修正。

【并行机制与性能解说】

• 线程/向量通道如何创建：这里没有 OS 线程。gang 是编译期/运行期概念，由 ISPC 编译器直接翻译成 SIMD 向量指令（--target=avx2-i32x8 → AVX2 的 8 个 32-bit lane）。8 个 “instance” 就是同一条向量指令的 8 个 lane，在同一个核上锁步推进。讲义明确提醒：前面这些代码只会跑在四核机器的一个核上；要多核，需要 ISPC 的另一个原语 task。
• 工作如何分配：foreach 的静态交错分配（静态 = 分配与数据无关，无动态负载均衡）。三轮循环的 i 值是 uniform 的，因此循环控制本身在向量化后只做一份。

• 共享数据如何处理：x 和 result 是 uniform float*（指针本身 uniform），但 x[i] 是 varying 表达式（每个 lane 不同）。例子中 lane 间完全不通信（每个输出元素只依赖自己的输入），所以不需要任何同步原语；如果要做归约，必须用跨 instance 通信原语 reduce_add——讲义用 sumall1/sumall2 的例子强调：sum 若是 uniform 而 x[i] 不是 uniform 表达式，会直接编译报类型错误（”Result: compile-time type error”）。

• Work / Span / 并行度分析（设单元素成本 c，含 terms 次内层迭代）：
  • Work（总工作量） T₁ = Θ(N · terms · c)。本例 N = 2²⁴、terms = 5。
  • 抽象的 Span（关键路径）：数据并行抽象里各迭代相互独立，故 T∞ = Θ(terms · c)（单个元素的成本）+ gang 返回的屏障延迟。抽象并行度 T₁/T∞ = Θ(N)，看起来”有 1600 万路并行”。
  • 实现的并行度：SIMD 实现下 gang 只有 programCount = 8 条真正并发的通道，且它们是锁步的：实际 T∞(实现) = Θ((N/8) · terms · c)，并行度(实现) = 8。
  • 结论：这就是”SMD 抽象 vs SIMD 实现”的量化含义——抽象告诉你”所有迭代相互独立、可任意调度”，实现只给你 8 路。 想拿到 4 核 × 8 路 = 32 路，必须用 ISPC task（或换用 OpenMP/CUDA 等模型），这正好呼应讲义对 gang 图的注释。
• 瓶颈：
  1. 访存模式（最大瓶颈）：交错版每条 packed load 搬 8×4 = 32 B（半条 cache line）且地址连续 → 高效；分块版各 lane 地址跨步 count（本例 count = N/8 = 2M 个 float = 8 MB 远）→ 需要 gather，讲义原话是”更复杂、更昂贵的指令”。实测差异来自内存事务数而不是算术量。
  2. 带宽（本例的真正天花板，见第 4 节算例）：N = 16M、每次迭代读 4 B 写 4 B，共 128 MB 流量；在 20 GB/s 的假设带宽下已经要 ≥6.4 ms，而算术只需约 0.65 ms（16 核 768 GFLOPS 的算例）——这道题是内存带宽题，不是浮点题。
  3. 条件控制流的掩码开销：本函数无分支，但若循环体里有 if，不同 lane 走不同路径时 SIMD 会串行化两条路径（谓词执行），有效 lane 利用率下降。
  4. sinx_blocked 里 count = N/programCount 依赖整除假设，N 不整除时尾部要单独处理（讲义用注释”assume N % programCount = 0”点明）。

3.2 示例 2：pthreads 分块求和 + 伪共享（共享地址空间模型）

// ============ 文件: sum_pthreads.cpp ============
// 编译（release）：g++ -O3 -march=native -pthread sum_pthreads.cpp -o sum_pthreads
// 运行：./sum_pthreads
#include <pthread.h>
#include <stdio.h>
#include <stdlib.h>

#define N        (1 << 26)          // 64M 个 float = 256 MB
#define NTHREAD  8
#define PAD      (64 - 8 - 8 - 8 - 8)   // 使结构体恰好 64 B = 1 条 cache line

// 每个线程的参数与结果：按 cache line 对齐并填充，避免 false sharing（伪共享）
typedef struct __attribute__((aligned(64))) {
    const float* x;
    long         begin;
    long         end;
    double       partial;       // 各线程私有部分和（不共享写）
    char         pad[PAD];
} arg_t;

static void* worker(void* p) {
    arg_t* a = (arg_t*)p;
    double s = 0.0;
    for (long i = a->begin; i < a->end; i++)   // 顺序遍历自己那一段
        s += a->x[i];
    a->partial = s;                            // 只写自己 cache line 内的字段
    return NULL;
}

int main(void) {
    float* x = (float*)aligned_alloc(64, (size_t)N * sizeof(float));
    for (long i = 0; i < N; i++) x[i] = 1.0f;  // 期望和 = 67108864

    pthread_t tid[NTHREAD];
    arg_t     arg[NTHREAD];

    long chunk = N / NTHREAD;                  // 分块（blocked）划分：每线程连续一段
    for (int t = 0; t < NTHREAD; t++) {
        arg[t].x     = x;
        arg[t].begin = t * chunk;
        arg[t].end   = (t == NTHREAD - 1) ? N : (t + 1) * chunk;
        arg[t].partial = 0.0;
        pthread_create(&tid[t], NULL, worker, &arg[t]);   // 创建 8 个线程
    }

    double total = 0.0;
    for (int t = 0; t < NTHREAD; t++) {
        pthread_join(tid[t], NULL);            // 等待（fork-join 的 join）
        total += arg[t].partial;               // 归约：串行加 8 个部分和
    }
    printf("total = %.0f (expected %d)\n", total, N);
    free(x);
    return 0;
}

【代码做什么？】

1. 分配 256 MB 的 x 并全部填 1.0。
2. 把数组分块切成 8 段（每段 32 MB 连续），每段配一个 64 字节对齐的 arg_t。
3. 主线程 pthread_create 创建 8 个 OS 线程；每个线程顺序遍历自己那段，累加到一个私有的 double partial。
4. pthread_join 隐式构成 fork-join 的 join 点；主线程把 8 个部分和相加。

【并行机制与性能解说】

• 线程如何创建：pthread_create 走 OS 系统调用 API，由 OS 把每个线程映射到硬件的执行上下文（execution context）（讲义最后一页思考题的答案：mapping 由 OS 负责）。8 个 pthreads > 4 核 × 2 上下文 = 8 个上下文时刚好占满；若开 16 个就会有过订阅，OS 需要分时切换（增加上下文切换与缓存污染）。
• 工作如何分配：分块（blocked）/ 静态划分。因为每线程拿到的是连续区间，硬件预取器与 TLB 都很友好——这与 ISPC 里”分块”的结论并不矛盾：那里的问题是分块破坏了 SIMD lane 的连续性，而这里每个线程是标量顺序执行，连续性由单线程自己的地址流保证。同一个”分块 vs 交错”的决策，在 SIMD 与多线程两种实现里结论不同——这正是”抽象 vs 实现”的最好例证。
• 共享数据如何处理：x 只读（所有线程读同一份，cache line 可在多个 L1/L2 里以 Shared 状态共存，无一致性流量）；partial 各线程私有，且用 64 字节填充隔离——如果 8 个 partial 挤在同一条 cache line 里，每个线程的写都会让该 line 在核间往返失效（cache line ping-pong），这就是伪共享（false sharing）：程序语义无误，性能却暴跌。
• Work / Span / 并行度：
  • Work：T₁ = Θ(N) 次浮点加 + N 次 4 B load。N = 2²⁶。
  • Span：T∞ = Θ(N/8)（一个线程的那一段）+ fork/join 开销（pthread_create 约 10 µs 量级 + pthread_join 同步）。注意：这不是 O(log N) —— 因为归约只在最后串行加 8 个数，O(8)；真正的关键路径是”一个线程干完自己那段”。
  • 并行度 = T₁/T∞ ≈ 8。但注意这个 8 是”上限”，不是”实际”：真正的约束是内存带宽（见下）。
  • 可扩展性上限：这是一个 bandwidth-bound 内核（每个元素只做 1 次加法，算术强度 1 flop / 4 B = 0.25 flop/byte）。加到 16 核不会更快，因为 DRAM 带宽不变。用 Amdahl 的话说，这里的”串行部分”是共享的 DRAM 带宽，不是代码里的某个临界区。
• 瓶颈（按重要性）：
  1. DRAM 带宽：读 256 MB（写只有 64 B），在本节假设的 20 GB/s 下 ≥12.8 ms。计算只需 64M 次加法 ÷ (8 核 × 3 GHz × 8 宽 SIMD) ≈ 0.33 ms。~40 倍差距 ⇒ 必须用向量化+多核把算力堆高是没用的，应该做的是减少访存（分块复用、融合）。
  2. 伪共享：若把 pad 去掉、8 个 double partial 相邻，则每次写都要独占整条 cache line；60+ ns 的 line 迁移 × 每线程只在末尾写一次……本例中每线程只写一次，影响小；但若循环体内累加的是共享的计数器（total += x[i] 且无锁），就会变成每次迭代一次 line 迁移，性能直接掉一到两个数量级（见 4.4 的数值算例）。
  3. 尾部负载不均：N % NTHREAD != 0 时最后一段略长，本例用 (t == NTHREAD-1) ? N : ... 处理；更普遍的不均（例如每段计算量不同）需要动态调度。
  4. 归约的串行尾巴：8 个部分和的相加是 O(NTHREAD) 串行，占比可忽略；但如果线程数上千（或在 GPU 上），就要用树形/对数步归约。
• 反例（讲义强调的”共享地址空间很容易搬起石头砸自己的脚”）：把上面改成”所有线程用一把互斥锁往同一个 total 里累加”，语义仍然正确，但每次加法都变成一次原子/加锁的串行临界区：此时 Span = Θ(N · t_critical)，并行度 ≈ 1，加核完全无效（4.5 有数值算例）。“程序正确但性能很差”是共享地址空间模型的典型失效模式。

3.3 示例 3：MPI 一维 stencil 的 halo 交换（消息传递模型）

// ============ 文件: stencil_mpi.c ============
// 编译（release）：mpicc -O3 -march=native stencil_mpi.c -o stencil_mpi
// 运行：mpirun -np 8 ./stencil_mpi
#include <mpi.h>
#include <stdio.h>
#include <stdlib.h>

#define WIDTH    1000000           // 全局网格宽度（每步需要邻居的 1 个元素）
#define NITER    100
#define TAG_L    100               // 消息 tag：来自左邻居
#define TAG_R    101               // 消息 tag：来自右邻居

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

    if (WIDTH % size != 0) { if (rank == 0) printf("WIDTH 必须被 size 整除\n");
                             MPI_Finalize(); return 1; }
    const int local = WIDTH / size;

    // cur/next 各含 2 个 halo 单元（下标 0 与 local+1）
    double* cur  = (double*)calloc(local + 2, sizeof(double));
    double* next = (double*)calloc(local + 2, sizeof(double));
    for (int i = 1; i <= local; i++) cur[i] = rank * 1000.0 + i;   // 初始化

    int left  = (rank - 1 + size) % size;      // 环状拓扑：边界绕回
    int right = (rank + 1) % size;

    MPI_Barrier(MPI_COMM_WORLD);
    double t0 = MPI_Wtime();

    for (int it = 0; it < NITER; it++) {
        // —— halo 交换：左右各一对 Sendrecv（可同时进行，避免死锁）——
        MPI_Sendrecv(&cur[1],       1, MPI_DOUBLE, left,  TAG_R,
                     &cur[local+1], 1, MPI_DOUBLE, right, TAG_R,
                     MPI_COMM_WORLD, MPI_STATUS_IGNORE);
        MPI_Sendrecv(&cur[local],   1, MPI_DOUBLE, right, TAG_L,
                     &cur[0],       1, MPI_DOUBLE, left,  TAG_L,
                     MPI_COMM_WORLD, MPI_STATUS_IGNORE);

        // —— 计算：5 点差分 cur[i] 的邻居在 cur[i±1] ——
        for (int i = 1; i <= local; i++)
            next[i] = 0.25 * (cur[i-1] + 2.0 * cur[i] + cur[i+1]);

        double* tmp = cur; cur = next; next = tmp;   // 交换缓冲区（零拷贝）
    }

    double t1 = MPI_Wtime();
    double local_sum = 0.0;
    for (int i = 1; i <= local; i++) local_sum += cur[i];
    double global_sum = 0.0;
    MPI_Allreduce(&local_sum, &global_sum, 1, MPI_DOUBLE, MPI_SUM, MPI_COMM_WORLD);

    if (rank == 0)
        printf("ranks=%d width=%d iters=%d  time=%.3f s  checksum=%.3f\n",
               size, WIDTH, NITER, t1 - t0, global_sum);

    free(cur); free(next);
    MPI_Finalize();
    return 0;
}

【代码做什么？】

1. MPI_Init 启动 size 个进程（不是线程），每个进程一个 rank；每进程负责全局网格的一段 local = WIDTH/size。
2. 每个进程的数组两端各留 1 个 halo（影子/晕圈）单元，用来存放邻居的边界值。
3. 每次迭代：(a) 与左右邻居各交换 1 个边界值（用 MPI_Sendrecv 一次调用同时完成收与发，避免”双方都先 recv 再 send”造成的死锁）；(b) 用 halo 更新自己内部的 local 个点；(c) 交换 cur/next 指针。
4. 结束时用 MPI_Allreduce 做一次全局归约求校验和。

【并行机制与性能解说】

• 进程/通信如何建立：mpirun -np 8 起 8 个独立进程，各有私有地址空间；它们之间没有共享内存可依赖（即使在同一节点上，MPI 也可能走共享内存通道，但编程模型上不可见）。数据交换只能靠 MPI_Sendrecv 显式拷贝。每对通信由 (发送者, tag) 匹配——本例用两个 tag 区分左/右邻居，避免消息错配。
• 工作如何分配：分块（blocked） 一维划分，每 rank 一段连续区间。这与共享地址空间的分块同理，但这里的”分块”必须在编译/启动期确定通信伙伴——这就是消息传递”结构强”的体现：读代码就能看出通信发生在哪里。
• 共享数据如何处理：无共享数据。cur/next 是各进程私有的；边界值通过消息复制过来。没有伪共享、没有缓存一致性流量，代价是每步都要显式拷贝 halo。
• Work / Span / 并行度：
  • Work：T₁ = Θ(NITER · (WIDTH 次差分计算))，即 100 × 10⁶ 次更新，每次约 4 flops ≈ 4×10⁸ flops。
  • Span（关键路径）：T∞ = Θ(NITER · (local · c_compute + α + halo_bytes/β))。”每个迭代都要等邻居的 halo” 形成逐步依赖链——这是 stencil 这类问题的关键路径本质：迭代之间不能并行，只有同一迭代内部的不同区间能并行。
  • 并行度 ≈ WIDTH（同一迭代内的点互相独立）。看起来极高，但实际加速被通信限制：T(p) ≈ T₁/p + NITER·(2α + 2h/β)，其中 h 是每个 halo 的字节数（本例只有 1 个 double = 8 B）。
  • 可扩展性上限（数值见 4.6）：halo 只有 8 B/次/邻居，却要付整个消息的固定延迟 α（约 2 µs），所以本例几乎完全是延迟受限：加 rank 会减少计算时间但通信延迟不变 ⇒ 存在一个最优 p。反过来，若 halo 面积/体积比很大（3D 大 stencil、或用很小的子域），计算/通信比会进一步恶化。
• 瓶颈（按重要性）：
  1. 通信延迟 α：8 B 的消息也要走完整个协议栈。增大每 rank 的计算量（更好的做法：增加每步的计算密度或用”波前/时间分块 temporal blocking”让 halo 复用多次）能摊薄它。
  2. 同步点：每次迭代两个 MPI_Sendrecv 都是同步点，MPI_Allreduce 更是全局同步；迭代间的依赖使流水线无法重叠。
  3. 负载不均：WIDTH % size != 0 直接报错退出；真实代码要么允许最后一段不等长，要么用一维/二维分解使各段均衡。
  4. 带宽 β（当 halo 变大时成为主因）：若每步交换的是”面”而不是”点”（3D 分解中常见），字节数按面积增长，此时 h/β 会主导。

3.4 示例 4：OpenMP 数据并行 map（数据并行模型 + 共享地址空间实现）

// ============ 文件: map_omp.cpp ============
// 编译（release）：g++ -O3 -march=native -fopenmp map_omp.cpp -o map_omp
// 运行：OMP_NUM_THREADS=8 ./map_omp
#include <omp.h>
#include <cstdio>
#include <cstdlib>
#include <cmath>

int main() {
    const long N = 1L << 24;                       // 16M 元素
    float* x = (float*)aligned_alloc(64, N * sizeof(float));
    float* y = (float*)aligned_alloc(64, N * sizeof(float));
    for (long i = 0; i < N; i++) x[i] = (float)((i % 2048) - 1024) / 1024.0f;

    double t0 = omp_get_wtime();

    // 数据并行 "map"：把 f 施加到集合 x 的每个元素上，写出到集合 y
    // schedule(static) = 静态分块（与 pthreads 示例同构的划分方式）
    #pragma omp parallel for schedule(static)
    for (long i = 0; i < N; i++) {
        float v = x[i];
        y[i] = (v < 0.0f) ? -v : v;                // f(v) = |v|
    }

    double t1 = omp_get_wtime();

    // 归约（数据并行的 reduce 模式；ISPC 里对应 reduce_add）
    double sum = 0.0;
    #pragma omp parallel for reduction(+:sum) schedule(static)
    for (long i = 0; i < N; i++) sum += y[i];

    double t2 = omp_get_wtime();
    printf("map: %.3f ms   reduce: %.3f ms   sum=%.0f\n",
           (t1 - t0) * 1e3, (t2 - t1) * 1e3, sum);

    free(x); free(y);
    return 0;
}

【代码做什么？】

1. 分配两个 64 MB 数组（共 128 MB），初始化 x。
2. 第一个 parallel for 是数据并行的 map：f(v)=\|v\| 施加到 x 的每个元素，结果写入 y——这正是讲义 slide 里 absolute_value 的 C++/OpenMP 版本（用 foreach 的 ISPC 版本语义完全一致）。
3. 第二个 parallel for 带 reduction(+:sum)：这是数据并行里的 reduce 模式；与 ISPC 的 reduce_add(partial) 对应——每个线程先在私有变量上累加，最后合并（讲义对 sumall2 的解释）。
4. 打印两段的耗时。

【并行机制与性能解说】

• 线程如何创建：OpenMP 采用共享地址空间模型 + fork-join：进入 parallel for 时”fork”出一个线程组（实现通常是线程池，不是每次真的 pthread_create），循环结束在隐式屏障处”join”。同步是隐含在 map 结束处的——与讲义对数据并行模型的定义完全一致（”Synchronization is implicit at the end of the map”）。
• 工作如何分配：schedule(static) 是静态分块（把迭代切成连续块分给线程，与示例 2 的 pthreads 分块同理）；schedule(dynamic) 则是”动态抢占式”分配，适合每迭代耗时差异大的场合，代价是每次取任务都有同步开销。
• 共享数据如何处理：x、y 共享但不重叠（每迭代写自己的 y[i]），无竞争；sum 由 reduction 子句自动变成”每线程私有副本 + 结束合并”，程序员看不到共享写。这里也顺带说明：“数据并行模型”在实现上几乎总是落在”共享地址空间的硬件”之上——讲义在总结页的原话是：数据并行”assumes a shared address space from which to load inputs/store results”。
• Work / Span / 并行度：
  • Map 的 Work = N 次”读 4 B + 比较/取负 + 写 4 B”；Span = 单元素成本 + 屏障延迟；并行度 = N（抽象层面）。
  • Reduce 的 Work = N 次加法；Span = 每线程的 N/p 次加法 + O(p) 的合并（若用树形归约则为 O(log p)）；并行度 = N/(N/p) = p——归约的并行度由线程数封顶，而不是由 N 封顶（因为同一线程内的累加是有依赖的串行链）。这是一个很容易答错的考点。
  • 可扩展性上限：Map 是纯流式访存（读 64 MB + 写 64 MB = 128 MB），完全受 DRAM 带宽限制：假设 20 GB/s，两段各需 ≥6.4 ms，与线程数无关；超过”带宽饱和点”后再加线程只会浪费。
• 瓶颈：
  1. DRAM 带宽（决定性）：两个 kernel 都几乎没有数据复用（算术强度分别约 0.25 和 0.125 flop/byte）。优化方向是融合（fusion）或分块（tiling）以减少字节数，而不是加线程。
  2. 隐式屏障：每次 parallel for 结束都有一个全部线程的屏障，线程数多、迭代短时屏障开销占比会上升（Barrier 成本约 1-10 µs 量级 + 不平衡等待）。
  3. 向量化：-O3 -march=native 下 y[i] = f(x[i]) 通常能向量化为 8 宽 AVX；若编译器因指针别名不敢向量化，可加 restrict 或 #pragma omp simd。这直接呼应示例 1：同一份”数据并行”语义，最终带宽受限于是否能生成 packed load/store。
  4. 动态调度的开销 / 负载不均：元素级 map 的迭代成本均匀，静态分块最省；若循环体里有分支或数据依赖的不均匀（稀疏结构），则需要 schedule(dynamic, chunk)。

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

本节所有”假设参数”都显式列出，属教学估算量级（非讲义原文数值）；讲义本身没有给出这些具体数字，请把它们当作可替换的建模输入。

4.1 work-span 模型与并行度

对并行程序的依赖图 DAG：

• Work T₁：用一个处理器执行全部工作所需时间（DAG 的总节点成本）。
• Span T∞（关键路径、critical path）：用无限多处理器执行所需时间（DAG 的最长路径成本）。
• 并行度（parallelism） = T₁ / T∞：能有效利用的处理器数上限。

• 调度界（greedy scheduler，如 work-stealing）：用 p 个处理器时

  T_p ≤ T₁/p + T∞        （且 T_p ≥ max(T₁/p, T∞)）
  

  ⇒ 当 p ≫ T₁/T∞ 时，T∞ 项主导，加处理器不再有用。这是”并行度上限”的定量表述。

表：本讲四个示例的 Work / Span / 并行度对照

4.2 Amdahl 定律与强/弱扩展

设可并行比例 f，串行比例 s = 1 - f，处理器数 p（忽略开销）：

  Speedup(p) = T₁ / T_p = 1 / ( s + (1 - s)/p )        lim_{p→∞} = 1/s

数值算例 A（Amdahl）：某程序有 2% 的串行部分（s = 0.02）。

• p = 8：Speedup = 1/(0.02 + 0.98/8) = 1/0.1425 = 7.02×（效率 87.7%）
• p = 16：Speedup = 1/(0.02 + 0.98/16) = 1/0.08125 = 12.31×（效率 76.9%）
• p = 64：Speedup = 1/(0.02 + 0.98/64) = 1/0.03531 = 28.32×（效率 44.3%）
• 上限：1/0.02 = 50×——无论多少核，都不可能超过 50 倍。

数值算例 B（锁的串行化，对应示例 2’）：8 个线程各做 10⁶ 次”加锁—自增—解锁”，每次临界区耗时 50 ns：

• 总临界区时间 = 8 × 10⁶ × 50 ns = 0.4 s（必须串行）
• 8 核并行本来只需 8×10⁶/8 × 1 ns ≈ 1 ms
• 实测 ≈ 0.4 s ⇒ 相对理想并行慢 ≈ 400 倍，加速比 ≈ 1.0
• 结论：临界区里哪怕只有几十纳秒，只要它在迭代里，Amdahl 的 s 就趋近 1。 共享地址空间模型的”性能陷阱”就是这种形状。

数值算例 C（弱扩展/Gustafson）：如果问题规模随 p 一起放大（每核工作量固定为 T），则 Speedup = s + (1-s)·p，p = 16、s = 0.02 时为 0.02 + 0.98×16 = 15.7×。这就是为什么超算都追求”弱扩展效率”——它绕过了 Amdahl 的上限，但要求问题本身能变大。

4.3 算术强度与 Roofline

算术强度（arithmetic intensity）：I = W / Q，其中 W 是浮点操作数，Q 是访问的字节数（DRAM 流量）。Roofline 模型：

  attainable_perf = min( P_peak , I × BW )
        ↑ 计算上限            ↑ 带宽上限（"屋顶"的斜边）
  拐点（ridge point）= P_peak / BW   [flops/byte]
     I < ridge  ⇒ memory bound（带宽受限，优化=减少访存/提高复用）
     I > ridge  ⇒ compute  bound（算力受限，优化=更宽 SIMD/更好指令混合）

数值算例 D（算力峰值）：讲义复习页给出的单芯片规格是”4 核 × 2 上下文 × 每时钟 2 条指令（其中 1 条 8 宽 SIMD）”。

  4 核 × 3.0 GHz × 8 宽 SIMD × 2 (FMA 计 2 flops) = 192 GFLOPS
  若放大到 16 核（同一时钟/宽度）：16 × 3.0e9 × 8 × 2 = 768 GFLOPS

数值算例 E（带宽 vs 一次遍历的成本，讲义反复使用的论证方式）：

  一个 256 MB 的数组（float32，64M 元素）
  假设 DRAM 带宽 = 20 GB/s
  只读一遍： 256 MB / 20 GB/s = 0.256e9 / 20e9 = 12.8 ms
  读+写一遍： 512 MB / 20 GB/s = 25.6 ms
  ⇒ 任何"整数组遍历一次"的并行算法，其时间下界就是这十几毫秒，
     与你有多少核、多少车道(SIMD) 无关。

数值算例 F（Roofline 判定，用示例 1 的 sinx）：设 terms = 5。

  每个元素的操作数（内层 5 次迭代）：
      value += sign*numer/denom   → 乘 + 加（除法按乘法倒数计）≈ 2 flops
      numer *= x*x                → 2 flops
      denom *= (2j+2)(2j+3)       → 2 flops（整数亦可）
      sign  *= -1                 → 可编译期展开，忽略
    小计 ≈ 6 flops/迭代 × 5 = 30 flops
    加上 x*x*x 的初始化 ≈ 2 flops        ⇒ W ≈ 32 flops/元素
  每个元素的 DRAM 流量：读 4 B + 写 4 B = 8 B   ⇒ I = 32/8 = 4.0 flops/byte
  拐点 = 768 GFLOPS / 20 GB/s = 38.4 flops/byte
  因为 4.0 ≪ 38.4 ⇒ 严重 memory bound
  可达性能上限 = I × BW = 4.0 × 20 GB/s = 80 GFLOP/s
              = 80/768 = 10.4% 的峰值算力
  ⇒ 结论：把 SIMD 从 8 宽加到 16 宽、把核数翻倍，对 sinx 几乎无益；
     真正的优化是(1)提高复用（分块，把同一元素供多个计算使用）、
     (2)减少流量（融合、用更低精度/更小数据类型）、(3)避免 gather。

表：常见内核的算术强度与受限类型（用上面的假设 BW = 20 GB/s、P_peak = 768 GFLOPS）

4.4 通信成本模型：延迟 α + 带宽 β（消息传递 & 缓存行迁移）

  Hockney 模型（点对点消息）：T_msg(n) = α + n/β
     α = 消息启动延迟（协议、软件栈、拷贝启动），β = 链路/拷贝带宽
  共享地址空间的"隐式通信"（缓存行迁移 / 伪共享）也有同一形式：
     T_line = α_cache + n_line/β_interconnect ， α_cache ≈ 60-100 ns 量级

数值算例 G（消息大小的影响）：α = 2 µs，β = 10 GB/s。

⇒ 小消息要么合并（把多次发送打包成一条），要么避免（提高计算/通信比）。这正是示例 3 中”halo 只有 8 B 却要付 2 µs”的定量解释。

数值算例 H（伪共享的代价）：4 个线程各自累加一个位于同一条 cache line 的 int 计数器，每线程 10⁷ 次。

  每次自增都需要"独占"该 cache line ⇒ 约 80 ns 的 line 迁移
  总时间 ≈ 4 × 10⁷ × 80 ns = 3.2 s
  填充后（每个计数器独占一条 line，line 常驻各自 L1 的 Exclusive 状态）：
     每线程 ≈ 10⁷ 次 × ~1 ns ≈ 10 ms      ⇒ 约 300 倍的差距
  ⇒ 这是"共享地址空间模型"里最典型、最难从代码上看出（但最容易用 padding 修好）的性能陷阱。

4.5 数值算例：示例 2 的端到端估算

假设：N = 2²⁶ = 67,108,864 个 float（256 MB），8 线程，DRAM 带宽 20 GB/s，8 核各 3.0 GHz、8 宽 SIMD、每时钟 1 条向量加。

  计算下界：67.1e6 次加法 / (8 核 × 3.0e9 × 8 宽) = 67.1e6 / 1.92e11 = 0.35 ms
  访存下界：256 MB / 20 GB/s = 12.8 ms
  ⇒ 预测耗时 ≈ 12.8 ms（内存受限，效率 = 0.35/12.8 = 2.7% 的算力利用率）
  若改为"只读一遍、只写一个部分和"的向量化实现，仍然 ≈12.8 ms（带宽就是墙）
  反例（全局锁版本）：8e6 次加锁 × 50 ns ≈ 0.4 s = 400 ms（比 12.8 ms 慢 31 倍）

4.6 数值算例：示例 3 的扩展性（寻找最优 rank 数）

假设：WIDTH = 10⁶ 个 double（8 MB），NITER = 100，单点更新约 4 flops，单核吞吐 3.0e9 × 8 宽 × 1 = 24 Gflop/s；α = 2 µs，halo = 8 B + 消息头（按 100 B 计），β = 10 GB/s。

  单 rank 计算时间 = 100 × 1e6 × 4 flops / 24e9 = 0.0167 s = 16.7 ms
  p = 8 时：计算 = 16.7/8 = 2.1 ms
            通信 = 100 次迭代 × 2 条消息 × (2 µs + 100 B/10 GB/s)
                 = 100 × 2 × 2.01 µs = 402 µs = 0.40 ms
            总计 ≈ 2.5 ms   （加速比 ≈ 6.7×）
  p = 64 时：计算 = 16.7/64 = 0.26 ms；通信仍 ≈ 0.40 ms
            总计 ≈ 0.66 ms  （相对 p=8 只快 3.8 倍，效率降到 24%）
  ⇒ 通信时间与 p 无关（每次迭代固定 2 条消息）⇒ 存在"收益递减"；
     p 增大到计算时间 ≈ 通信时间（约 0.4 ms ⇒ p ≈ 42）附近即最优。
  改进方向（对应讲义"abstraction distance / 通信成本要反映现实"的思想）：
     (1) 增大每 rank 的数据量（弱扩展）；(2) 时间分块让 halo 复用多次；
     (3) 用一次大消息换掉多次小消息（合并 halo）；(4) 非阻塞通信与计算重叠。

5. 关键要点

1. 永远区分”编程模型（抽象）”与”硬件实现”。ISPC 的 gang 抽象 = 8 条逻辑指令流，实现 = 一条 8 宽 AVX 指令；pthreads 的线程抽象，实现 = OS 把线程映射到硬件执行上下文。性能藏在实现的细节里，不要把抽象语义当成性能保证（讲义的原话：”conflating abstraction with implementation is a common cause for confusion”）。
2. 三种通信抽象各用结构换性能可预测性：共享地址空间结构最少、最自然，但所有 load/store 在代码上看起来一样贵，实际却天差地别（局部/远端、有无伪共享、是否争锁）；消息传递把通信全部显式化成消息，结构约束反而帮你更快写出可扩展的正确程序；数据并行把计算钉成”map over collection”，用刚性换自动并行与编译期优化（预取、融合）。
3. 抽象必须”低到能反映通信成本，又高到能保持可移植”（abstraction distance 原则）。这是讲义总结页的核心判据，也是理解为什么现代实践是混合模型：节点内用共享地址空间（pthreads/OpenMP），节点间用消息传递（MPI），kernel 内再加受限的共享内存同步（CUDA/OpenCL）。
4. 硬件尺度决定有效模型：单核内 SIMD 宽度决定数据并行的粒度（gather 与 packed load 的差别就是 10 倍级访存成本）；片内多核共享 L3/DRAM 带宽决定”加核是否有用”；多路节点是 NUMA（局部性重要）；集群只能消息传递。讲义用 Intel i7 的 ring、Niagara 2 的 crossbar、Altix UV1000 的 fat tree（4096 核单一地址空间）、El Capitan 的 dragonfly + MI300A 统一地址空间，说明”支持共享地址空间的硬件代价随规模急剧上升”。
5. 性能上限先算带宽与并行度，再谈优化：对每个内核先估 I = W/Q 与 T∞。若 I < P_peak/BW（本例 sinx：4.0 < 38.4），则优化方向是减少字节数（融合/分块/更小数据类型）而非加宽 SIMD；若 T₁/T∞ 很小（如全局锁版本 ≈ 1），则优化方向是缩短关键路径（消除串行化）而非加处理器。

6. 常见陷阱与注意事项

• 把”抽象语义允许并行”当成”实现能并行”：foreach 声明 N 个迭代相互独立，但 ISPC gang 只有 programCount（= SIMD 宽度）路；要在多核上跑必须用 ISPC 的 task 或换用多线程模型。同理，#pragma omp parallel for 写出来了也不代表生成了向量指令——没向量化就等于丢掉了 8 倍。
• 数据竞争（data race）与”看起来正确”的共享内存程序：共享地址空间是顺序编程的自然延伸，但两个线程对同一变量无同步的并发读写是未定义行为（讲义用 while (x == 0) {} 轮询的例子开场）。更隐蔽的是“程序正确但很慢”：比如用一把全局锁把并行循环串行化（示例 2’，实测比理想并行慢数百倍），代码 review 往往查不出来。
• 伪共享（false sharing）：不同线程写不同变量但落在同一条 cache line 上，导致 line 在核间反复迁移。语义完全正确，性能可能差两个数量级。修法是按 cache line（典型 64 B）对齐填充，或让每线程使用私有副本 + 末尾归约。注意它只发生在”共享地址空间 + 缓存一致”的实现里；消息传递模型没有这个问题，但代价是显式拷贝。
• 忽视带宽与算术强度：大量”逐元素”内核（map、求和、DAXPY、stencil）的算术强度都在 1 flop/byte 以下，性能由 DRAM 带宽封顶，与核数/SIMD 宽度无关（示例 E：256 MB 数组单次遍历 ≥12.8 ms）。在这类内核上做”把 SIMD 加宽、把线程数翻倍”的优化是纯粹的浪费；正确方向是提高复用（分块/时间分块）或减少流量（kernel 融合、压缩数据类型）。
• 负载不均与静态划分的盲区：静态分块在每迭代成本均匀时最省（无调度开销），但一旦成本不均（稀疏结构、分支、边界条件、尾部不整除），最慢的那个线程决定总时间（T_p ≥ max_i T_i）。反之，动态调度有每任务一次的同步开销，任务粒度过细时会退化（对应 schedule(static) vs schedule(dynamic, chunk) 的选择）。
• 消息传递里的死锁与错配：两个 rank 都先做阻塞 MPI_Recv 再 Send 会死锁（示例 3 用 MPI_Sendrecv 规避）；tag 使用不当会让消息匹配到错误的接收者（尤其在有多个通信阶段时）；小消息频繁发送会被固定延迟 α 主导（算例 G：8 B 消息的有效带宽只有 0.004 GB/s）。先想通信结构，再写代码——这正是消息传递模型”强制结构”的价值所在。
• 非确定的数据并行写法：讲义用 shift_negative 说明——多个迭代可能写同一位置，而数据并行模型不规定迭代顺序、也不提供细粒度互斥，因此该程序是非确定的。在 foreach/kernel 里写”跨迭代依赖”（如邻接写、累加到共享位置）必须改用其他模型或显式原子操作，代码可能”有时对”，但这不代表它是对的。

7. 思考题（带答案）

问题 1. 一个只在 ISPC 里用 foreach 写成的 sin(x) 计算程序，在 4 核 8 宽 SIMD 的机器上实测速度只有理论峰值的 ~3%。请分别从”抽象/实现”、”访存模式”、”带宽”三个角度解释原因，并给出至少两条可行的优化方向。

【答案】 (1) 抽象 vs 实现：foreach 的抽象语义是”N 个迭代彼此独立、可任意调度”，看起来并行度是 N；但 ISPC gang 的实现只是一条 8 宽 SIMD 向量指令流，跑在一个核上，实际并行度只有 programCount = 8。要真正吃满 4 核，必须用 ISPC 的 task 原语（或改写成 OpenMP/CUDA）。也就是说，实际并行度是 8 而不是 4×8 = 32，仅此一项就损失了 4 倍。 (2) 访存模式：如果按讲义给的手写”分块（blocked）”版本写（start = programIndex * count），同一轮各个 lane 访问的地址相隔 count 个元素、并不连续，编译器必须生成 gather（_mm_i32gather，2013 年才随 AVX2 出现，AVX2 甚至没有 SIMD scatter）而不是一条 packed load。gather 的访存事务数成倍增加、延迟更高，有效带宽远低于连续向量访存。改成 foreach/交错访问后访存变成连续 32 B packed load，这一步就能拿到数倍提升。 (3) 带宽与算术强度：terms = 5 时每元素约 32 flops，而每个元素只读 4 B、写 4 B，故 I ≈ 4.0 flops/byte；而该机器的 Roofline 拐点 = 768 GFLOPS / 20 GB/s = 38.4 flops/byte。I ≪ 拐点 ⇒ 严重 memory bound，性能上限 = I × BW = 80 GFLOP/s ≈ 10.4% 峰值；加上 gather 与未向量化等因素，实测 3% 完全在预期内。 可行优化：(a) 把索引改成交错（foreach）以获得 packed load/store；(b) 分块复用数据、把多个元素的计算做成一次访存喂多个计算（提高 I）；(c) 减少内层运算（把 denom/sign 用 uniform 变量、把除法换成乘以事先算好的 1/denom，必要时用多项式/查表）；(d) 用 float 而非 double 降低字节数；(e) 若允许近似，换成更少 terms 的多项式并 -ffast-math 让编译器向量化。

问题 2. 你的程序在一台 8 核共享地址空间机器上做 y[i] = f(x[i])（N = 2²⁴ 个 float），8 线程实测 6.4 ms；把线程数加到 16 后仍然 6.4 ms。请判断这是什么瓶颈，用数字论证，并说明若把问题换成”每元素做 500 次浮点运算”会有什么不同。

【答案】 这是 DRAM 带宽瓶颈，不是并行度不足。数字论证：N = 2²⁴ = 16.78M 个 float，读写各 4 B ⇒ 总流量 = 2 × 67.1 MB = 134.2 MB。实测 6.4 ms ⇒ 有效带宽 = 134.2 MB / 6.4 ms = 21 GB/s，正好等于该机器的 DRAM 带宽（本节假设 20 GB/s 量级）。带宽是共享资源，加线程不会增加它，所以 8→16 线程毫秒不动；同时说明 8 线程时就已经把带宽打满了。 若每元素改为 500 flops：I = 500/8 = 62.5 flops/byte > 拐点 38.4 ⇒ 变成 compute bound。此时 8 线程的计算时间 = 16.78e6 × 500 / (8 核 × 3.0e9 × 8 宽 × 2 FMA) = 8.39e9 / 384e9 ≈ 21.8 ms，远大于带宽下界 6.4 ms ⇒ 时间由算力决定，加线程（到 16）会接近线性加速（约 11 ms），直到再次撞上带宽或达到 16 核的算力上限。结论：同一个”数据并行 map”结构，因算术强度不同而落在 Roofline 的两侧，优化手段完全不同——先算 I 再动手。

问题 3. 讲义说”编程模型与机器类型之间的对应关系是模糊的”。请举出两个方向的例子，并解释”abstraction distance（抽象距离）”为什么既不能太大也不能太小。另外：为什么现代超算普遍采用”节点内共享地址空间 + 节点间消息传递”的混合模型？

【答案】 两个方向的例子：(a) 在共享地址空间的硬件上实现消息传递：极其常见。此时”发送消息”就是从消息库缓冲区拷贝内存，”接收消息”就是从库缓冲区拷回——通信仍然成立，只是走的是本地拷贝而不是网络。(b) 在硬件不支持共享地址空间的机器上实现共享地址空间抽象：用软件（如共享虚拟内存/软件 DSM）——把含有共享变量的页标记为 invalid，由缺页异常处理程序（page-fault handler） 发出网络请求把页面取到本地，程序仍然能”读写共享变量”，只是每次缺页都要经历一次极慢的软件路径。 abstraction distance 的两难：抽象距离太小（抽象几乎照抄某个硬件的细节，比如把”某个具体 SIMD 宽度、某个具体 cache line 大小”写进模型），性能可预测、容易调优，但代码不可移植，换个机器就废；抽象距离太大（模型对所有访存/通信一视同仁，例如”所有 load/store 一样贵”），代码可移植、好写，但程序员无法从源码判断性能，优化只能靠猜。讲义的原则是：保持足够低以让性能可预测，足够高以保留灵活性与可移植性——共享地址空间模型正是”距离偏大”的典型（好写但容易写出慢程序），消息传递模型则”距离偏小”（难写但可扩展）。 为什么超算是混合模型：硬件上，同一节点内（多核 + 共享 L3 + 缓存一致）能以极低成本提供共享地址空间，所以用 pthreads/OpenMP 最划算、编程最舒服；而跨节点若要在硬件上维持全局共享地址空间，成本极高（讲义：SMP”统一地差”、NUMA 已经需要程序员处理局部性；Altix UV1000 用 fat tree 换来 4096 核单一地址空间，代价是巨额互连开销），因此节点间只提供消息通信，用 MPI 显式传递，让程序员面对真实的通信成本。于是”节点内共享内存 + 节点间消息传递”成了现代实践（讲义原话：”very, very common in practice”），这也是 CUDA/OpenCL 里”kernel 之间用数据并行、同一 core/block 内的线程用共享内存通信”这一安排的同一个道理——在每个尺度上选用与硬件成本最匹配的抽象。