Lecture 14: Performance Analysis / Profiling

目录 · ← l13 · l15 →

Lecture 14: Performance Analysis / Profiling

1. 章节标题与概述

Lecture 14: Performance Analysis / Profiling（性能分析与剖析）

• 本讲核心问题：怎么知道程序慢在哪里，以及怎么证明自己知道？ 并行程序的性能问题有四大类来源——计算（compute）、内存带宽/延迟（memory bandwidth / latency）、同步（synchronization）、并行度不足/负载不均（insufficient parallelism / load imbalance）；这四类问题的现象常常互相伪装（”加线程没变快”既可能是带宽饱和，也可能是锁竞争）。本讲要建立一套可测量、可证伪的方法学：先用正确的计时与统计得到可信的数字，再用剖析（profiling）与硬件性能计数器（PMU）定位热点与瓶颈类型，最后用高水位（high watermark）实验与 Roofline 模型把”瓶颈诊断”从直觉变成定量结论。

• 涉及的主要硬件/软件机制：硬件侧是 PMU（Performance Monitoring Unit，性能监控单元）——处理器内部一组可编程的事件计数器（退休指令数、时钟周期、L1/L2/LLC 命中与缺失、从内存控制器读取的字节数、各类 stall 周期），以及计数器溢出中断这一把”计数器”变成”采样剖析器”的机制；软件侧是插桩（instrumentation）、统计采样（sampling）、计数器读取（counting / event-based）三大家族工具（gprof、perf、Intel VTune、PAPI、PCM、Callgrind、Nsight Compute/Systems），以及把测量结果转成优化决策的分析框架（Roofline、Amdahl、work-span、strong/weak scaling）。

• 在并行计算知识体系中的角色：本讲是全课程的方法论枢纽。前面的讲座（指令级并行、SIMD、多核、缓存一致性、伪共享、同步）提供了”性能为什么会这样”的机制，后面的讲座（性能优化、异构与 GPU、专用化）提供”应该怎么改”的手段；本讲回答的是中间那个问题——“我手上这份代码，实际卡在哪一条机制上？” 它把三个此前零散出现的判据合并成一条工作流：①用 work-span / 并行度判断”是不是并行度不够”；②用算术强度 / Roofline 判断”是不是带宽不够”；③用 PMU 计数器 + 高水位实验判断”是不是同步或延迟不够”。此外，本讲也是后续”自动性能优化”（Halide autoscheduler、自动调优、LLM 代码生成 + profile 反馈）得以成立的前提：所有自动优化循环的本质都是”测量 → 建模 → 搜索 → 再测量”，而没有可信的测量，搜索就失去目标函数。

• 配套材料：

  • CMU 15-418 Fall 2026 本讲讲义：未发布。 Fall 2026 已公开的讲义 PDF 位于 https://www.cs.cmu.edu/~418/lectures/ 之下（可直接下载）；但 Lecture 14（Performance Analysis / Profiling）在 Fall 2026 尚未在该公开目录发布，课程主页 https://www.cs.cmu.edu/~418/ 与日程表 https://www.cs.cmu.edu/~418/schedule.html 中本讲日期为 Sep 28，其材料栏为空。
  • 历史学期对应讲义（09_perfeval.pdf / 13_perftools.pdf，F22 归档）：需登录、未公开。 它们位于 /afs/cs/academic/class/15418-*/public/ 之下，需要 CMU 登录；本仓库 supp/ 目录下的 f22_perfeval.pdf、f22_perftools.pdf 两个文件只是 CMU IdP 登录页的 HTML（各约 1.7 KB），不含任何讲义内容，因此本笔记无法引用其正文。
  • 关于本讲可用的抽取文本的说明：cs149_supp/thoughtprocess.txt 与 cs149_supp/aiperfoptimization.txt 是姊妹课程 CS149 的相邻主题讲义（”并行编程的思维过程”与”领域特定系统与自动性能优化”），并不是 CMU 15-418 Lecture 14 的讲义正文。本笔记以其中可直接复用的术语、公式、案例与结论（Amdahl 定律、分解/分配/编排、图像处理的工作计数与算术强度结论、高水位与 Roofline 方法学、PMU 计数器、profile 驱动的优化循环）作为事实基础，其余关于剖析工具机制（gprof 的 mcount 插桩、perf 的采样与 perf_event 子系统、PMU 多路复用、Nsight Compute 指标、Callgrind 模拟）的论述基于该主题的成熟公开知识（如 Roofline 原始论文 Williams et al. 2009 与各工具官方文档所描述的机制）；所有定量结论均显式给出假设参数，便于读者自行复核。
  • 已公开的姊妹课程（Stanford CS149）材料（本笔记的主要事实基础）：
    • cs149_supp/perfopt2.txt：已公开。CS149 Fall 2025 Lecture 6 “Performance Optimization Part II: Locality, Communication, and Contention”（68 页）。它包含本讲最核心的三块方法论：“确定你受限于计算、内存带宽还是同步”、高水位实验、Roofline 模型（computing-limited 水平段 vs bandwidth-limited 斜段），以及 PMU 计数器（Intel PCM 的 C++ API 示例：getIPC、getL3CacheHitRatio、getBytesReadFromMC）与 VTune / PAPI / oprofile 的指引；并附有固定问题规模测速的陷阱（32 处理器 SGI Origin 2000 上的 258×258 网格）与超线性加速的解释。
    • cs149_supp/thoughtprocess.txt：已公开。CS149 Fall 2025 Lecture 4 “Parallelizing Code: The Programming Thought Process”。提供分解/分配/编排/映射四阶段框架、Amdahl 定律（S = 串行比例 ⇒ 最大加速比 ≤ 1/S）、静态 vs 动态分配、ISPC foreach/gang 抽象、交错（interleaved，packed load）vs 分块（blocked，gather）分配、网格求解器与屏障，以及”每元素加锁 → 每次迭代加锁一次“与”三个屏障 → 一个屏障“的经典优化。
    • cs149_supp/aiperfoptimization.txt：已公开。CS149 Fall 2025 Lecture 13 “Domain-Specific Programming Systems and Automatic Performance Optimization”。提供图像处理的精确定量工作计数（3×3 二维卷积 9×W×H；两趟分离 6×W×H；分块后 (34/16)×3×W×H = 6.4×W×H）、“两趟实现比二维实现算术强度低 2 倍”这一结论、分块/融合的生产者-消费者局部性论证、Halide 的算法/调度分离与自动调度器（代价模型为小型 MLP，输出 27 个系数供手写代价模型使用；1.4M 个 schedule 在 166 秒内评估完）、以及用 profiling 统计量驱动的 LLM 反思式优化循环（示例统计量：Timing: 32 ms、SM util 42%、DRAM util 89%、L2 cache hit 68%）。
    • cs149_supp/efficiency.txt：已公开。提供加速比与效率的定义（speedup = T(1)/T(P)；效率 = 每单位资源的性能，如每瓦/每美元/每芯片面积的 ops/s）与”FAST ≠ EFFICIENT（在 10 核上只有 2× 算不算好结果？）”的判据。
  • CMU 侧其它已公开的相关讲义（用于术语与课程语境对齐）：extracted/07_progperf1.txt、extracted/08_progperf2.txt（Performance Optimization I/II，含周期性自我剖析与重分配）、extracted/05-CUDA-programming.txt（GPU 剖析语境）。
  • 未发布/需登录：Fall 2026 的录像（Panopto/YouTube）在日程表中被注释隐藏，属未发布；Ed 讨论区、Autolab、Canvas 均需登录。考试覆盖方面，公开考试页显示 Exam 2 覆盖 Lectures 14–24（不含客座讲座），本讲是其起点。
  • 说明：讲义首页有时写有历史学期（如 Fall 2025）字样，这是讲义沿用，属正常现象，不视为错误。

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

2.1 性能的第一性度量：到底在测什么时间

• 定义与目的：性能分析的第一步是选对被测对象。三个不同的”时间”回答三个不同的问题：
  • wall-clock time（墙钟时间 / 真实经过时间）：从开始到结束的物理时间。这是用户唯一关心的时间，也是加速比公式里必须用的时间。std::chrono::steady_clock、MPI_Wtime、omp_get_wtime 测的都是它。
  • CPU time（进程占用 CPU 的时间）：所有线程/进程的时间之和。clock()、getrusage()、/usr/bin/time 的 user+sys 给的是它。陷阱：在 8 线程并行区里，clock() 会报告约 8× 的墙钟时间——在并行程序里拿 clock() 算加速比，会得到荒谬的”负加速比”。
  • 硬件事件计数（event counts）：退休指令数、周期数、cache miss 数、DRAM 字节数。它们不是时间，但能解释时间：T = 指令数 / IPC / 频率，T ≥ 流量 / 带宽。

• 直观解释（”它是什么？”）：把程序想成一家餐厅。墙钟时间是顾客从进门到出门的时长（顾客只在乎这个）；CPU 时间是所有厨师在厨房里干活的总人时（8 个厨师一起干 1 分钟的菜，总人时是 8 分钟，但顾客只等了 1 分钟）；硬件计数器是厨房里的各种仪表——切菜次数、冰箱开门次数、走到仓库取货的次数。要诊断”为什么顾客等这么久”，你必须同时看墙钟（等太久）、看人时（厨师在忙还是互相撞）、看仪表（是不是都在排队等同一个冰箱）。

• 性能特征（延迟 / 带宽 / 吞吐量）：这三者在本讲中始终要分清。延迟（latency）是单次操作的时间（DRAM ~80–100 ns）；带宽（bandwidth）是单位时间能搬多少数据（DRAM ~20 GB/s 量级）；吞吐量（throughput）是单位时间完成多少任务。延迟靠”隐藏”（并发、流水、预取），带宽只能靠”少搬”——这是贯穿全课程的区分。

2.2 剖析方法学三大家族：插桩、采样、计数器

• 定义与目的：剖析器（profiler）回答”时间花在哪个函数/哪行代码/哪条指令上”。三大家族的机制截然不同，代价与可信度也不同。

• 直观解释（”它是什么？”）：
  • 插桩（instrumentation）＝在每个房间装一台打卡机：进函数打一次卡、出函数打一次卡，于是精确知道每段代码的耗时与调用次数。代价是打卡机本身要时间（每次函数调用多几十到上百纳秒），而且它会改变被测量的对象（尤其对小函数和深递归）。
  • 采样（sampling）＝总经理每隔 1 毫秒随机巡楼一次，记下”此刻谁在干活”。它不改变程序（代价 <1%），但得到的是统计估计而非精确值；样本越多越准，短命函数可能一次都没被看到。
  • 计数器计数（counting / event-based）＝看水电表：不关心”谁在干活”，只关心”总共用了多少度电、开了多少次冰箱”。它给出精确的总量（周期数、缺失数、DRAM 字节数），但不能直接告诉你这些量属于哪一行代码——除非配合采样（把”哪个事件溢出”作为采样触发条件），这正是 perf record -e LLC-load-misses 的做法。
• 架构/机制图解（工具链数据通路，硬件结构视角）：

        ┌──────────────────────────── 一个 CPU Core ────────────────────────────┐
        │                                                                        │
        │   ┌──────────────┐   取指/译码   ┌───────────────────────────────┐      │
        │   │ Frontend     │──────────────▶│ Backend: 执行单元 / LSU / FPU │      │
        │   └──────┬───────┘               └───────────┬───────────────────┘      │
        │          │ 事件脉冲                            │ 事件脉冲                 │
        │          │ (retired inst, branch-miss,        │ (L1 miss, LLC miss,     │
        │          │  frontend stall...)                │  stall-cycles...)       │
        │   ┌──────▼────────────────────────────────────▼─────────────┐           │
        │   │  PMU：事件选择寄存器 + N 个通用计数器 + 固定计数器          │           │
        │   │  (可编程: perf_event_open() / MSR 写入; 溢出阈值可设)      │           │
        │   └──────────────────────────┬───────────────────────────────┘           │
        └──────────────────────────────┼───────────────────────────────────────────┘
                                       │ ① 计数器溢出 → NMI 中断（采样模式）
                                       │ ② 直接读计数器（计数模式）
        ┌──────────────────────────────▼───────────────────────────────────┐
        │  内核：perf_event 子系统（进程/线程绑定、多路复用 multiplexing、     │
        │        上下文切换时保存/恢复计数器、权限控制）                       │
        └──────────────────────────────┬───────────────────────────────────┘
                                       ▼
        ┌──────────────────────────────────────────────────────────────────┐
        │  用户态工具层                                                      │
        │  perf stat / perf record+report / Intel VTune / PAPI / PCM /      │
        │  likwid / oprofile / NVIDIA Nsight Compute(设备侧 SM 计数器)        │
        └──────────────────────────────┬───────────────────────────────────┘
                                       ▼
        输出：IPC、cache 命中率、DRAM 字节数、stall 分解、调用栈火焰图、roofline 点

• 关键操作与性能特征：
  • 计数器数量是稀缺资源：一个物理 core 通常只有 4–8 个通用计数器，而可命名的事件有几百个。想同时测的事件超过物理计数器数时，内核会做多路复用（multiplexing）——把时间切片分给不同事件组，再按启用比例外推。后果：计数结果带外推误差（通常几 %，负载不平稳时更大），perf stat 会打印 [63.45%] 这样的”该事件实际被启用时间的占比”，这一列必须看。
  • 采样的开销几乎恒定，与函数调用频率无关；插桩的开销与调用次数成正比。对调用 10⁸ 次的小函数，插桩版本可能慢 10 倍以上。
  • 采样有”滑移（skid）”：中断是流水线延迟之后才被处理和记录指令指针的，因此被归因的指令可能比真正引起事件的指令晚几十到上百条。对单条热点指令的定位要小心，对函数级归因通常无伤。
• 表格：三类剖析技术对比

2.3 硬件性能计数器（PMU）：事件选择与诊断语义

• 定义与目的：PMU 是把”微架构内部发生了什么”暴露给软件的唯一官方通道。它解决的问题是：在不改代码、不猜的前提下，拿到”程序在硬件上到底干了什么”的客观证据。CS149 的 perfopt2 讲义对此的表述是：”所有现代处理器都有低层事件性能计数器——记录指令完成数、时钟周期、L2/L3 缓存命中/缺失、从内存控制器读取的字节数等；例如 Intel 的 Performance Counter Monitor Tool 提供访问这些寄存器的 C++ API（getIPC、getL3CacheHitRatio、getBytesReadFromMC）；另见 Intel VTune、PAPI、oprofile”。

• 直观解释（”它是什么？”）：PMU 就是汽车的仪表盘。你不能通过”感觉车很肉”来修发动机，但你能看转速、水温、涡轮压力、油耗（每百公里多少升）。”油耗高 + 转速低”指向一个诊断，”转速高 + 车速低”指向完全不同的诊断。单个计数器几乎从不下结论，一组计数器的比值才下结论——这正是 IPC、cache 命中率、stall 分解、字节/秒这些”派生指标”存在的意义。

• 表格：关键 PMU 事件与它们诊断什么

• 架构/机制图解（采样剖析器的时间轴与归因，软件执行模型视角）：

  采样剖析器工作原理：以频率 F (如 1 kHz, 即每 1 ms) 触发中断，记录"此刻正在执行谁"

  时间轴 (ms)  0     1     2     3     4     5     6     7     8     9    10
               |-----|-----|-----|-----|-----|-----|-----|-----|-----|-----|
  正在执行     [main ][k_A  ][k_A  ][k_B  ][k_A  ][sync ][k_B  ][k_A  ][k_B ][k_A ]
                ▲     ▲     ▲     ▲     ▲     ▲     ▲     ▲     ▲     ▲
             10 个样本，每个样本 = 一次"栈回溯 + 归因"

  自顶向下归因（火焰图，宽度正比于样本数）
  ┌───────────────────────────────────────────────────────────────────┐
  │                              main                                │  10/10
  ├───────────────────────────────────────┬───────────────────────────┤
  │              kernel_A                 │        kernel_B           │  5/10 | 3/10
  ├──────────────┬────────────┬───────────┼──────────────┬────────────┤
  │  loop_over_i │  loop_over_j│  tail    │   dev_1      │   dev_2    │
  └──────────────┴────────────┴───────────┴──────────────┴────────────┘
                                    │
                           1/10 是 sync（同步/屏障/锁/等待）  ← 这一类占比是并行程序最关心的
  自底向上（调用者视角）：谁"发起"了这些时间

• 统计灵敏度的定量后果（必须记住的算术）：采样数的统计误差按二项分布估计，绝对误差 m = z·sqrt(p(1−p)/n)。取 z = 1.96（95% 置信）、p = 0.5（最坏情形）：

  想要的绝对精度	需要的样本数 n	1 kHz 下需要的测量时长
  ±5 个百分点	≈ 385	≈ 0.39 s
  ±1 个百分点	≈ 9604	≈ 9.6 s
  ±0.1 个百分点	≈ 960,000	≈ 960 s（16 分钟）

  推论 ①：“这个函数占 0.3% 时间”这种结论，在 1 秒的采样里根本不可信（0.3% × 1000 样本 = 3 个样本，泊松相对误差 ≈ 1/√3 ≈ 58%）。推论 ②：想用采样精确定位一个只占 0.1% 的函数，需要约 10⁴–10⁵ 个样本、即 10–100 秒的测量，或者改测更大/更重的输入。推论 ③：先用长跑粗筛、再用小输入精确定位热点行，是标准做法。

2.4 Roofline 与算术强度：把”瓶颈类型”变成坐标系

• 定义与目的：算术强度（arithmetic intensity） AI = 浮点操作数 / 从内存系统搬运的字节数（FLOP/Byte）。Roofline 模型（Williams et al. 2009，CS149 讲义直接引用该图）给出可达性能上界：

  可达性能 = min( 峰值算力 P_peak , AI × 可持续带宽 BW )
  ridge point（拐点）= P_peak / BW
  

  它解决的问题是：在动任何优化之前，先判断”加算力有用还是减流量有用”。AI < ridge ⇒ 带宽受限，任何算术优化的收益都接近 0；AI > ridge ⇒ 计算受限，减少访存无益，要提 ILP/SIMD/减少指令数。

• 直观解释（”它是什么？”）：Roofline 是一张“水管与漏斗”图。屋顶的水平段是处理器的算力上限（漏斗口径固定：每秒最多漏 768 GFLOP）；斜段是内存带宽的斜坡（水管流量固定：每秒最多供 20 GB 数据）。每搬来 1 字节数据，你能做多少次运算决定了你站在斜坡上还是屋顶下：算术强度太低，你搬来的数据”不够算”，处理器在等数据（斜坡限制）；算术强度够高，数据供得上，你受限于算力（屋顶限制）。

• 架构/机制图解（Roofline，双对数坐标）：

  可达性能
  (GFLOP/s,  fp32)
    768 ┤━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━━  屋顶：compute bound
        │                                          ┌──────  (16核×3.0GHz×8宽×2 FMA)
        │                                        ╱
        │                                      ╱   斜坡斜率 = 可持续带宽 = 20 GB/s
    100 ┤                                    ╱
        │                                  ╱
        │                                ╱
     30 ┤● 3×3 二维 blur (AI=2.25)     ╱
        │ ● 融合/分块两趟 (AI=1.50)   ╱
     15 ┤  ● 朴素两趟 blur (AI=0.75) ╱
        │    ● SAXPY (AI=0.083)    ╱
      1 ┤      × 伪共享自增 (AI≈0.016)
        └────┬────────┬──────────┬───┬──────────────────────────────────────▶
            0.083    0.75      2.25 38.4 = 768/20          AI (FLOP/Byte)
                                    ▲
                             ridge point：越过它才进入 compute bound

  优化方向的语言：
    斜段上的点 → 唯一有效的手段是"每字节做更多运算"：分块(tiling)、融合(fusion)、
                 提高数据复用、缩短数据宽度(fp32→fp16)、降低精度需求
    接近屋顶点 → 手段变成：SIMD 宽度、FMA、减少指令数、提高 ILP、消除除法/超越函数

• 性能特征：Roofline 的上界是乐观的（它假设带宽可以被完全利用、且不区分延迟与带宽），因此它给出的是“最好可能有多好”，即所谓 high watermark（高水位）——这正是它最大的用途：当你已经达到 Roofline 上界的 80%–90% 时，就应该停止微调，去做别的；当你离上界很远，模型会告诉你差距是”流量”还是”算力”造成的。

2.5 高水位方法学：把”我觉得”换成”我测到”

• 定义与目的：来自 CS149 perfopt2 讲义的核心方法论（原话：“永远、永远、永远先写最简单的并行版本，然后测量性能，看看你站在哪里”、“确定你的性能是受限于计算、内存带宽（或内存延迟），还是同步”、“尝试建立高水位：实践中最好能做到多少？你的实现离最好情形有多近？”）。高水位 = 通过故意破坏某一种资源的使用方式，测出”若这种资源不再受限，程序能跑多快”。

• 直观解释（”它是什么？”）：像给病人做激发试验。想判断”病人跑不动是不是因为心脏”？让他吸纯氧再跑一次——如果速度没变，那瓶颈就不是供氧。高水位实验就是给程序做激发试验：故意把访存全部改成命中同一个地址（等于把内存流量降到接近 0），如果程序快了很多 ⇒ 流量是瓶颈；如果几乎没变 ⇒ 流量不是瓶颈。

• 表格：四类高水位实验（讲义给出的四个，原文照录并解释）

• 讲义给出的最重要的一条免责声明：“计算、访存与同步几乎从不会完美重叠；因此总体性能很少完全由 compute 或 bandwidth 或 sync 之一决定。即便如此，性能对上列改动的敏感度仍是主导成本的很好指示。” 这句话定义了本讲的方法论诚实性：性能分析给的是归因与敏感度，不是”单一原因的判决书”。

2.6 可扩展性度量：加速比的正确测法（speedup measurement）与三类经典错误

• 定义与目的：
  • 加速比（speedup）：S(P) = T(1) / T(P)（T 必须是墙钟时间；讲义亦给出等价形式”P 个处理器上的 ops/s ÷ 单处理器 ops/s”）。加速比测量（speedup measurement）是本讲的核心技能之一——它看起来只是一次除法，却是性能报告中最常出错的地方（见本节末的三类错误）。
  • 效率（efficiency）：E(P) = S(P)/P——”性能 per 单位资源”，讲义强调资源可以是芯片面积、美元、瓦特。
  • strong scaling（强可扩展）：问题规模固定，看时间如何随 P 下降；weak scaling（弱可扩展）：每处理器的问题规模固定，看时间如何随机器变大保持不变。
  • Amdahl 定律：设 f 为本质上串行的工作占比，则 S(P) ≤ 1/(f + (1−f)/P)，且 S(∞) = 1/f。讲义中的直观版本：S = 串行比例 ⇒ 最大加速比 ≤ 1/S。
  • Gustafson 定律（弱扩展）：S_scaled(P) = P − f·(P−1)（随机器变大而变大问题规模时，串行部分只占固定绝对时间）。

• 直观解释（”它是什么？”）：Amdahl 是”买更大的厨房，做同一桌菜”——如果那道必须由主厨一个人完成的摆盘要 10 分钟，你请 100 个副厨也只能把其余部分压到接近 0，总时间仍 ≥ 10 分钟，最大加速比被钉死在 1/f。Gustafson 是”买更大的厨房，做更大的宴席”——串行部分（比如主厨尝味道）的绝对时间不变，但总工作量随厨师数线性增长，于是规模化的加速比可以接近 P。这解释了一个看似矛盾的现象：超级计算机的”实测加速比”常常好于 Amdahl 的强扩展预测，因为人们实际上在解更大的问题。

• 架构/机制图解（work-span / fork-join：判断”并行度够不够”）：

  Work-Span 模型：把 DAG 的"总节点数"与"最长依赖链"分开看

         ┌──── t0 (fork) ────┐
         │        │          │
   ┌─────▼──┐ ┌───▼────┐ ┌───▼────┐ ┌────────┐     4 个"任务"
   │ task0  │ │ task1  │ │ task2  │ │ task3  │
   │ w = 3  │ │ w = 3  │ │ w = 3  │ │ w = 3  │     每任务工作量 3
   └─────┬──┘ └───┬────┘ └───┬────┘ └────┬───┘
         │        │          │           │
         └────────┴────┬─────┴───────────┘
                  ┌────▼────┐
                  │ join    │  ← 屏障：Span 里必须计入
                  │ w = 1   │
                  └────┬────┘
                       ▼
   Work W = 3+3+3+3+1 = 13        Span S = 3 + 1 = 4
   并行度 (parallelism) = W/S = 3.25    ← 即使给 100 个核，最优时间也只降到 S

   上界（贪心调度定理）： T(P) ≤ W/P + S      ⇒  S(P) ≤ W/(W/P + S)
   "够不够并行"的判据：W/S ≫ P 才值得加核；W/S 与 P 同量级时已到天花板

• 三类经典的”测加速比”错误（来自 CS149 perfopt2 讲义，原文均有对应）：
  1. 基线选择错误：“常见陷阱：把并行程序的加速比，对比于’并行算法在单核上运行’。” 讲义原文：Common pitfall: compare parallel program speedup to parallel algorithm running on one core (easier to make yourself look good). 正确做法是与最好的（best known）串行实现比较，并在报告中写明基线是什么。注意：并行算法常常要做更多总工作（例如红黑着色法比 Gauss-Seidel 收敛更慢），把它单核的时间当基线会凭空”造出”加速比。
  2. 固定问题规模：32 处理器 SGI Origin 2000 上跑 258×258 网格（每处理器仅约 310 个格点）时，讲义明确标注 “No benefit! (slight slowdown)”——通信/计算比太高，问题对机器来说太小；同一台机器上换成 1K×1K 网格（每处理器 ~32K 格点）才看到正常扩展。推论：用固定问题规模评价机器是危险的，因为规模一变，负载均衡、开销、算术强度、局部性全都变。
  3. 超线性加速（super-linear speedup）的误读：S(P) > P 通常不是”并行效率超过 100%”，而是工作集变小导致的：随着 P 增大，每个处理器分到的数据块开始装进 cache（讲义原文：”with enough processors, chunk of grid assigned to each processor begins to fit in cache”），或原本在单机上换页到磁盘（thrashing）的工作集在更大内存的机器上不再换页。它不是算法变好，而是被比较的两个系统在存储层次上不是同一个系统。

2.7 抽象与实现（abstraction pitfalls）：为什么”同一段抽象代码”性能差 3 倍

• 定义与目的：本讲的关键词之一是 abstraction pitfalls（抽象的陷阱）。抽象（abstraction）给程序员一个简洁的心智模型，但性能由实现（implementation）决定，而实现常常泄漏到抽象之外。CS149 的 thoughtprocess 讲义用 ISPC 给出了教科书级的例子：同一份 sinx 的抽象语义，两种”分配（assignment）”策略编译出的机器码完全不同：
  • 交错分配（interleaved）：idx = i + programIndex，8 个 program instance 访问连续的 8 个 float ⇒ 一条 packed vector load（vmovaps，对应 _mm256_load_ps） 即可完成；
  • 分块分配（blocked）：start = programIndex*count; idx = start + i，8 个 instance 访问相隔 count 的元素，在内存中不连续 ⇒ 必须用 gather 指令（vgatherdps，对应 _mm256_i32gather_ps），讲义原文：”gather is a more complex, and more costly SIMD instruction…“。
  • 这就是”抽象 vs 实现”的完整教训：foreach 抽象只说”这些迭代要做”，没规定谁做哪些；而 ISPC 现实的实现对 foreach 采用的是静态交错方案（讲义原文：”abstraction leaves room for dynamic assignment, but current ISPC implementation is a static scheme”）。程序员若不知道实现，就无法解释自己程序的性能。

• 直观解释（”它是什么？”）：抽象像餐厅菜单上的”招牌炒饭”。菜单不写米是当季新米还是陈米、灶是几个火眼，但你吃到的东西由后厨决定。同一个”炒饭”抽象，两家店能差三倍。性能分析的重要一半工作，就是掀开后厨看一眼（看汇编、看 profiler 的指令级归因、看 perf stat 的 IPC 与缺失率）。

• 其它常见的抽象陷阱（本讲会反复出现）：
  • std::vector<T>::operator[] / 迭代器与越界检查：debug 构建与 release 构建的抽象完全不同（_GLIBCXX_ASSERTIONS）。
  • -O3 下的内联：被内联的函数在采样剖析的调用栈中消失，其时间被归到调用者头上。若确实需要区分，用 __attribute__((noinline)) 或看带 DWARF 的 perf report --inline。
  • 编译器优化掉你要测的代码：C++ 里未经使用的结果会被整段删除（dead code elimination），于是你测出”0 ms”的惊人性能。必须把结果”消费掉”（打印校验和、写入 volatile 全局、asm volatile("" ::"r"(x))）。
  • 计时器抽象：clock() 是 CPU 时间而 std::chrono::steady_clock 是墙钟；在多线程里用错会把加速比符号都搞反。
  • omp_get_wtime vs MPI_Wtime vs CUDA event：不同计时域的默认行为与同步语义不同（CUDA 的 cudaEvent 需要在正确 stream 上记录方有异步语义）。
  • 性能计数器命名的抽象：cache-misses 在不同微架构上映射到不同事件；跨机器比较原始计数没有意义，只有比值（IPC、命中率、字节/秒、% of peak）才可比较。

3. 代码示例与性能分析

3.1 示例 1：可复现的计时骨架 + 两趟 vs 分块融合 blur 的 Roofline 定位（C++ / OpenMP）

本示例严格对应 CS149 aiperfoptimization 讲义中的图像处理案例：3×3 盒式模糊（box blur），二维卷积 ↔ 两趟可分离实现 ↔ 分块融合实现，并给出每一版的工作量计数、流量估计、Roofline 定位。

// blur_bench.cpp —— 两趟(朴素) vs 分块融合 3x3 盒式模糊 + 可信计时骨架
//
// 编译(必须 release 优化): g++ -O3 -march=native -fopenmp blur_bench.cpp -o blur_bench
// 运行:                    OMP_NUM_THREADS=8 ./blur_bench 1024 20
// 剖析(PMU 计数):          perf stat -e cycles,instructions,L1-dcache-load-misses,LLC-load-misses ./blur_bench 1024 20
// 采样热点:                perf record -g -F 999 ./blur_bench 1024 20 && perf report

#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <cmath>
#include <vector>
#include <algorithm>
#include <chrono>
#include <omp.h>

static const float W3 = 1.0f / 3.0f;    // 1x3 与 3x1 的归一化权重

// ---------- 版本 A：两趟（先水平、再垂直），需要完整 (H+2) x W 的中间缓冲 ----------
// 对应讲义 slide "Two-pass 3x3 blur"
static void blur_x(const std::vector<float>& in, std::vector<float>& tmp, int W, int H) {
    const int stride = W + 2;
    #pragma omp parallel for schedule(static)
    for (int j = 0; j < H + 2; ++j) {
        const float* row = &in[(size_t)j * stride];
        float*       out = &tmp[(size_t)j * W];
        for (int i = 0; i < W; ++i)
            out[i] = W3 * (row[i] + row[i + 1] + row[i + 2]);
    }
}

static void blur_y(const std::vector<float>& tmp, std::vector<float>& out, int W, int H) {
    #pragma omp parallel for schedule(static)
    for (int j = 0; j < H; ++j) {
        const float* r0 = &tmp[(size_t)(j + 0) * W];
        const float* r1 = &tmp[(size_t)(j + 1) * W];
        const float* r2 = &tmp[(size_t)(j + 2) * W];
        float*       o  = &out[(size_t)j * W];
        for (int i = 0; i < W; ++i)
            o[i] = W3 * (r0[i] + r1[i] + r2[i]);
    }
}

static void blur_twopass(const std::vector<float>& in, std::vector<float>& out,
                         std::vector<float>& tmp, int W, int H) {
    blur_x(in, tmp, W, H);
    blur_y(tmp, out, W, H);
}

// ---------- 版本 B：分块融合（chunked / fused）----------
// 对应讲义 slide "Two-pass image blur, chunked (version 2)"：
// 只为 CHUNK+2 行输入生产中间结果，立刻消费掉，中间缓冲常驻 cache。
static void blur_chunked(const std::vector<float>& in, std::vector<float>& out,
                         int W, int H, int CHUNK) {
    const int stride  = W + 2;
    const int nchunks = (H + CHUNK - 1) / CHUNK;

    #pragma omp parallel
    {
        std::vector<float> tb((size_t)(CHUNK + 2) * W);   // 线程私有中间缓冲（关键：私有！）
        #pragma omp for schedule(dynamic, 1)
        for (int c = 0; c < nchunks; ++c) {
            const int j0    = c * CHUNK;
            const int rows  = std::min(CHUNK, H - j0);    // 本块要产出的输出行数
            const int trows = rows + 2;                   // 因此需要的中间行数

            // 生产：中间结果行 j0 .. j0+trows-1
            for (int jj = 0; jj < trows; ++jj) {
                const float* row = &in[(size_t)(j0 + jj) * stride];
                float*       o   = &tb[(size_t)jj * W];
                for (int i = 0; i < W; ++i)
                    o[i] = W3 * (row[i] + row[i + 1] + row[i + 2]);
            }
            // 消费：输出行 j0 .. j0+rows-1（此时中间结果还在 cache 里）
            for (int jj = 0; jj < rows; ++jj) {
                const float* r0 = &tb[(size_t)(jj + 0) * W];
                const float* r1 = &tb[(size_t)(jj + 1) * W];
                const float* r2 = &tb[(size_t)(jj + 2) * W];
                float*       o  = &out[(size_t)(j0 + jj) * W];
                for (int i = 0; i < W; ++i)
                    o[i] = W3 * (r0[i] + r1[i] + r2[i]);
            }
        }
    }
}

// ---------- 计时骨架：warmup + 多次重复 + 中位数 + 防死代码消除 ----------
static double now_s() {
    using clk = std::chrono::steady_clock;                       // 墙钟，非 CPU 时间
    return std::chrono::duration<double>(clk::now().time_since_epoch()).count();
}

int main(int argc, char** argv) {
    const int W   = (argc > 1) ? std::atoi(argv[1]) : 1024;
    const int H   = W;
    const int REP = (argc > 2) ? std::atoi(argv[2]) : 20;
    const int CHUNK = 16;                                        // 讲义示例用 CHUNK_SIZE = 16

    std::vector<float> in((size_t)(W + 2) * (H + 2));
    std::vector<float> tmp((size_t)(H + 2) * W);
    std::vector<float> out((size_t)W * H, 0.0f);
    for (size_t k = 0; k < in.size(); ++k)
        in[k] = (float)((k * 2654435761u) % 1000) * 0.001f;

    // 双精度累加校验和：既验证正确性，又阻止编译器删除整个计算（防 DCE）
    auto checksum = [](const std::vector<float>& v) {
        double s = 0.0;
        for (float x : v) s += (double)x;
        return s;
    };

    auto bench = [&](const char* name, auto&& fn) {
        fn();                                                    // warmup：页错误/首次触碰/TLB
        std::vector<double> t;
        for (int r = 0; r < REP; ++r) {
            const double t0 = now_s();
            fn();
            t.push_back(now_s() - t0);
        }
        std::sort(t.begin(), t.end());
        printf("%-14s median = %8.3f ms   min = %8.3f ms   max = %8.3f ms\n",
               name, t[REP / 2] * 1e3, t.front() * 1e3, t.back() * 1e3);
        return t[REP / 2];
    };

    printf("W=H=%d  threads=%d  CHUNK=%d\n", W, H, omp_get_max_threads(), CHUNK);
    const double t2p = bench("two-pass", [&]{ blur_twopass(in, out, tmp, W, H); });
    const double ck1 = checksum(out);
    out.assign((size_t)W * H, 0.0f);
    const double tch = bench("chunked",  [&]{ blur_chunked(in, out, W, H, CHUNK); });
    const double ck2 = checksum(out);

    // ---------- 把时间翻译成硬件语言 ----------
    const double pixels = (double)W * H;
    const double macs_ideal = 6.0 * pixels;                      // 讲义计数：6 × W × H (MAC)
    const double macs_chunk = 3.0 * (CHUNK + 2.0) / CHUNK * pixels
                            + 3.0 * pixels;                      // 讲义：((CHUNK+2)+CHUNK)/CHUNK × 3 × W × H

    // DRAM 流量模型（B）：输入读一次 + 输出写一次 = 固有流量
    const double B_intrinsic = (double)(W + 2) * (H + 2) * 4 + pixels * 4;
    // 朴素两趟：中间缓冲被完整写出又被完整读回（远超 cache 容量）
    const double B_twopass  = B_intrinsic + 2.0 * (double)(H + 2) * W * 4;

    printf("\n[流量与算术强度]\n");
    printf("  intrinsic traffic = %.2f MB   two-pass traffic = %.2f MB (比值 %.2fx)\n",
           B_intrinsic * 1e-6, B_twopass * 1e-6, B_twopass / B_intrinsic);
    printf("  AI(two-pass) = %.3f FLOP/B   AI(chunked) = %.3f FLOP/B\n",
           2 * macs_ideal / B_twopass, 2 * macs_chunk / B_intrinsic);

    printf("\n[实测]\n");
    printf("  two-pass: %7.3f ms -> %7.2f GFLOP/s, %6.2f GB/s\n",
           t2p * 1e3, 2 * macs_ideal / t2p / 1e9, B_twopass / t2p / 1e9);
    printf("  chunked : %7.3f ms -> %7.2f GFLOP/s, %6.2f GB/s\n",
           tch * 1e3, 2 * macs_chunk / tch / 1e9, B_intrinsic / tch / 1e9);
    printf("  speedup(chunked / two-pass) = %.2fx\n", t2p / tch);
    printf("  checksum: %.3f vs %.3f  (diff = %.3e)\n", ck1, ck2, std::fabs(ck1 - ck2));
    return 0;
}

• 【代码做什么？】
  1. 分配输入 in（(W+2)×(H+2)，比输出多一圈，避免内层循环做边界判断）、中间缓冲 tmp（(H+2)×W）、输出 out（W×H）。
  2. 版本 A（blur_twopass）：先对所有行做水平 1×3 卷积写入 tmp，再对所有输出行做垂直 3×1 卷积写入 out。这是讲义 “Two-pass 3x3 blur” 的直接实现，中间缓冲大小 W×(H+2)。
  3. 版本 B（blur_chunked）：把输出按 CHUNK 行分块。对每一块，只生产 rows+2 行中间结果（放在线程私有的 tb 里，大小只有 (CHUNK+2)×W ≈ 74 KB，可常驻 L2），立刻消费掉产出 rows 行输出。这是讲义 “chunked (version 2)”：用分块把生产者-消费者局部性捕获在 cache 里，”trends to ideal value as CHUNK is increased”。
  4. 计时骨架：steady_clock 测墙钟；先 warmup() 一次（处理页错误、首次触碰、首次分配）；重复 REP 次，打印中位数与实际 min/max（min 反映”无干扰”性能，max 反映抖动，中位数反映典型值——三者都看，才知道测量是否可信）；每次迭代都对同一份输出做双精度校验和并打印，既验证两版结果一致，又防止编译器把整个 kernel 删掉。
  5. 最后把时间翻译成硬件语言：打印 MB、FLOP/B、GFLOP/s、GB/s 与两版加速比。
• 【并行机制与性能解说】
  • 线程与工作分配：两个版本都用 #pragma omp parallel for 把行（或块）分配给线程。blur_x/blur_y 用 schedule(static)（每线程连续若干行，工作量均匀，无需动态调度开销，且保证了行内连续访问）；blur_chunked 用 schedule(dynamic,1) 是因为块的成本虽相同但 cache 行为抖动大，动态调度可以吸收抖动（代价是每次取任务一次原子操作，块数 H/CHUNK = 64，总开销可忽略）。
  • 共享数据如何处理：所有线程只读 in、只写自己负责的 out 区域（行/块不重叠），因此没有任何数据竞争，也不需要锁。blur_chunked 的中间缓冲 tb 被声明在 #pragma omp parallel 区域内 ⇒ 每个线程一份私有副本，这是该版本正确性的关键；若把它放到 parallel 区之外共享，就会产生数据竞争。
  • 向量化：内层 for (i...) 是对连续内存的无依赖循环，-O3 -march=native 会把它自动向量化为 8 宽 AVX2 的 vfmadd/vaddps。注意这条循环的访存模式：out[i] = W3*(row[i]+row[i+1]+row[i+2]) 对固定行是连续访问 ⇒ 编译器能用 packed load（这正是 thoughtprocess 讲义 slide 10 强调的”交错分配 ⇒ 一条 vmovaps“）；而如果改成”每个线程负责一组列“，就会退化成 gather。
  • Work / Span / 并行度：
    • Work：W = Θ(W·H) 次乘加（常数因子：两趟 6 MAC/像素，分块 6(1+1/CHUNK) MAC/像素）。1024×1024 时约 6.29 MFLOP（两趟）与 6.68 MFLOP（CHUNK=16）。
    • Span：关键是这里有几个同步点。blur_x 与 blur_y 之间存在一个隐式的全局依赖（垂直遍历需要所有行的水平结果），因此 naive 两趟版本的 Span 至少是”一趟的水平 Span + 一趟的垂直 Span”。单趟的 Span：每个线程处理 rows ≈ H/P 行、每行 W 次连续迭代，若强制线程内串行 ⇒ Θ(H·W/P)；再加上 parallel for 结束时的屏障（集中式 Θ(P)，树形 Θ(log P)）。所以 Span_two_pass ≈ Θ(H·W/P + P)
    • 并行度 = Work/Span = Θ(W·H) / Θ(H·W/P + P)。取 W=H=1024, P=16：Span 项 1024×1024/16 = 65536, P = 16 ⇒ 并行度 ≈ 1.05e6/6.55e4 ≈ 16——与线程数同量级，说明这份实现的并行度刚好”够用但不富余”；一旦 P 继续增大，H·W/P 变小而 P（屏障）不变甚至增大，并行度不再增长，这就是讲义 thoughtprocess slide 35 里 T = N²/P + P 那条曲线的来源（下一节会给出完整数值）。
    • 分块版本的 Span 更短：融合后两趟之间的全局依赖被打碎成逐块的局部依赖，块之间彼此独立（dynamic 调度下甚至可乱序完成），Span ≈ Θ((CHUNK+2)·W + P)，与 H 无关。
  • 瓶颈分析：W=H=1024 时，按 20 GB/s 与 768 GFLOP/s 的机器模型：
    • 两趟：流量 16.80 MB ⇒ 带宽下界 840 µs；算力时间 12.58 MFLOP/768 GFLOP/s ≈ 16 µs ⇒ 带宽受限约 51 倍。
    • 分块：流量回落到 8.40 MB ⇒ 带宽下界 420 µs ⇒ 预期加速比约 2.0×（讲义对两趟实现的算术强度结论”2x lower arithmetic intensity“就是这一条）。
    • 立刻可验证的诊断：程序打印的 GB/s 若已接近本机实测的可持续带宽（用 STREAM 类内核测），说明带宽已饱和，再加线程无用；若远低于可持续带宽，则瓶颈在延迟/访存并行度（MLP）或同步，而不是带宽。用 perf stat -e LLC-load-misses,stalled-cycles-backend 可以区分这两者（见 §2.3 表）。

3.2 示例 2：用 PMU 计数器诊断”加线程反而更慢”——伪共享的真实共享（pthreads + OpenMP）

// false_sharing.cpp —— packed(伪共享) vs padded(独占 cache line) 两种计数器布局
//
// 编译(必须 release): g++ -O3 -march=native -pthread -fopenmp false_sharing.cpp -o fs
// 运行:               ./fs packed 4 16777216
//                     ./fs padded 4 16777216
// 剖析:               perf stat -e cycles,instructions,cache-misses,LLC-load-misses ./fs packed 4 16777216
//                     perf c2c record ./fs packed 4 16777216 && perf c2c report   # 直接指出伪共享行

#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <thread>
#include <vector>
#include <atomic>
#include <chrono>
#include <algorithm>

struct alignas(64) PaddedCounter { volatile long long v; char pad[64 - sizeof(long long)]; };

static double now_s() {
    using clk = std::chrono::steady_clock;
    return std::chrono::duration<double>(clk::now().time_since_epoch()).count();
}

int main(int argc, char** argv) {
    const bool packed = (argc > 1) && (strcmp(argv[1], "packed") == 0);
    const int  T      = (argc > 2) ? std::atoi(argv[2]) : 4;
    const long long ITER = (argc > 3) ? std::atoll(argv[3]) : 16777216LL;

    // packed: T 个 long long 挤在 1-2 条 cache line 里（65 字节以下就落在同一行）
    // padded: 每个计数器独占 64 B
    std::vector<long long>  packed_cnt(T, 0);
    std::vector<PaddedCounter> pad_cnt(T);

    auto worker_packed = [&](int t) {
        volatile long long* p = &packed_cnt[t];
        for (long long i = 0; i < ITER; ++i) (*p)++;    // volatile 防止被优化成一次加法
    };
    auto worker_padded = [&](int t) {
        volatile long long* p = &pad_cnt[t].v;
        for (long long i = 0; i < ITER; ++i) (*p)++;
    };

    const double t0 = now_s();
    {
        std::vector<std::thread> th;
        for (int t = 0; t < T; ++t) {
            if (packed) th.emplace_back(worker_packed, t);   // 两个 lambda 类型不同，
            else        th.emplace_back(worker_padded, t);   // 不能写成三目表达式
        }
        for (auto& x : th) x.join();
    }
    const double dt = now_s() - t0;

    long long total = 0;
    for (int t = 0; t < T; ++t) total += packed ? packed_cnt[t] : pad_cnt[t].v;

    printf("%-7s threads=%2d  iter/thread=%lld  time = %8.3f s   aggregate = %8.2f Mops/s\n",
           packed ? "packed" : "padded", T, ITER, dt, (double)(T * ITER) / dt / 1e6);
    printf("total = %lld  (预期 %lld)\n", total, (long long)T * ITER);
    return 0;
}

• 【代码做什么？】
  1. T 个线程各自执行 ITER 次自增；packed 模式下所有计数器是 std::vector<long long> 里的连续元素——当 T ≤ 8 时它们全部落在同一条 64 B cache line 内；padded 模式用 alignas(64) + 填充让每个计数器独占一条 cache line。
  2. 用 volatile 阻止编译器把循环优化成一条 += ITER（这是这个微基准最关键的正确性细节）。
  3. 测墙钟、打印总吞吐量（Mops/s）与总和正确性校验。
• 【并行机制与性能解说】
  • 硬件层面发生了什么：x86 的一致性协议（MSI/MESI/MOESI，本课程 Lecture 11–13 主题）以 cache line（64 B） 为所有权粒度。packed 模式下，线程 0 要写 cnt[0]，必须让包含该字节的整条 line 处于 M（Modified，独占可写） 态；此刻线程 1 想写 cnt[1]——同一个 line 的另一个字节——就必须先把 line 从线程 0 的 cache 中无效化（I）并搬到自己的 cache。于是每次自增都触发一次line 迁移（cache-to-cache transfer + 失效 + 所有权串行化），代价约 几十到约 100 ns，并且把两个线程彻底串行化。

  • 架构/机制图解（伪共享的 line 乒乓，硬件结构视角）：

      共享的一条 64 B cache line（4 个计数器 + 邻居数据）
      ┌───────────────────────────────────────────────────────────────┐
      │ cnt0 │ cnt1 │ cnt2 │ cnt3 │  (同一行的其它变量 …)                │
      └───────────────────────────────────────────────────────────────┘
         ▲                                       ▲
         │ 线程0 要写 cnt0                        │ 线程1 要写 cnt1
      ┌──┴───────┐    ① 请求独占(M)            ┌─┴────────┐
      │  Core 0  │ ───────────────────────────▶ │  Core 1  │
      │ L1: M    │ ◀─── ② line 迁移 + Core1 变 I │ L1: I    │
      └──────────┘                               └──────────┘
            ③ Core1 再次请求 M ──▶ Core0 变 I ──▶ line 又搬回去 …… 循环往复
      结果：每次自增 = 一次跨核 line 迁移（~ 数十至 100 ns）
            吞吐量被钉在 ≈ 1 / 迁移延迟 ≈ 1e7 次/秒 量级，且 与线程数无关（不升反降）
    
      修复后（padded）：每个计数器独占一行
      ┌───────────────┐ ┌───────────────┐ ┌───────────────┐ ┌───────────────┐
      │ cnt0 + 60B pad│ │ cnt1 + 60B pad│ │ cnt2 + 60B pad│ │ cnt3 + 60B pad│
      └───────────────┘ └───────────────┘ └───────────────┘ └───────────────┘
         Core0: M            Core1: M           Core2: M          Core3: M
      各自命中自己的 L1 行，无一致性流量 ⇒ 吞吐量回到 ALU/内存极限（高 1~2 个数量级）
    

  • Work / Span / 并行度：
    • Work W = T × ITER 次自增（本程序里每个自增是 1 次读-改-写，视为 1 单位工作）。
    • Span：padded 模式下各线程互不干扰，Span = 单个线程的工作 = Θ(ITER) ⇒ 并行度 = W/S = T（完美的线性可扩展）。packed 模式下，每次自增都可能是一次 line 迁移，其延迟必须串行化，因此 S_packed = Θ(T × ITER × L_transfer)（L_transfer ≈ 数十~100 ns），并行度 = W/S ≈ 1/L_transfer（即约 10⁷ 次/秒，与 T 无关）——这就是”并行度从 T 掉到 1 以下“的严格含义：加线程不会变快，只会让 line 搬家更频繁。
  • 数量级估算（用于理解，实测请运行代码）：T=4, ITER=2²⁴（每线程 1677 万次）：
    • packed：总自增 6.71×10⁷ 次，若每次平均 ~100 ns 的迁移延迟 ⇒ ≈ 6.7 s；
    • padded：6.71×10⁷ 次 / (4 线程 × 3.0 GHz × ~1 次/周期) ⇒ ≈ 5.6 ms（外加每个计数器每 64 B 只有一份 line，无一致性流量）；
    • 比值约 10²–10³ 量级。注意实际数值强烈依赖微架构（有些实现会用”写独占但可延迟”的优化），所以必须实测——这正是本讲的立场。
  • 瓶颈类型：这不是带宽瓶颈（数据本来就在 cache 里，DRAM 字节数很低），而是一致性/同步导致的串行化瓶颈。用计数器区分二者的方法：伪共享的特征是 LLC-load-misses 低但 cache-misses/一致性事务高、且时间随线程数增加而变差；真实带宽瓶颈的特征是 从内存控制器读到的字节数已接近机器上限。perf c2c（cache-to-cache）子命令就是为此设计的：它直接给出发生了 line 争用的地址。
  • 正确修法（比 padding 更好）：不要共享写，要”私有写 + 归并”——这正是讲义 thoughtprocess slide 68→69 的教训：把”每个 (i,j) 迭代都加锁”改成”先累加到线程本地 myDiff，每次迭代只加锁一次“，原文标注 “Now only lock once per thread, not once per (i,j) loop iteration!”。在本例中对应 #pragma omp parallel for reduction(+:total) 或每线程局部累加后一次原子归并。一般化的判据：三个条件同时成立才需要动手——①多个线程写、②这些写在同一条 line 上、③这些写频率很高；只满足前两条（例如只在初始化/收尾各写一次）则成本可忽略。

3.3 示例 3：用设备侧计数器驱动优化（CUDA + Nsight Compute 反馈循环）

本示例把 aiperfoptimization 讲义中的”profile → 反思 → 改代码 → 再 profile“循环落到具体代码与指标上（讲义给出的反馈回路示例统计量是：Timing: 32 ms、SM util 42%、DRAM util 89%、L2 cache hit rate 68%）。

// saxpy_prof.cu —— 同一个 SAXPY 的两个版本：标量 + 网格跨步 vs float4 向量化 + __restrict__
//
// 编译: nvcc -O3 -arch=sm_80 -lineinfo saxpy_prof.cu -o saxpy
// 运行: ./saxpy 16777216 20
// 剖析(关键三件套: SM 吞吐、DRAM 吞吐、L2 命中率):
//   ncu --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed,\
//dram__throughput.avg.pct_of_peak_sustained_elapsed,\
//lts__t_sector_hit_rate.pct  ./saxpy 16777216 1
// 也可以用: ncu --set full ./saxpy 16777216 1     (更慢，但一次给全)

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

// 版本 A：一线程一元素（标量、4 B/线程/流）
__global__ void saxpy_scalar(int n, float a, const float* __restrict__ x, float* __restrict__ y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) y[i] = a * x[i] + y[i];
}

// 版本 B：网格跨步 + float4（16 B/线程/流，一次事务取满 128 B sector）
__global__ void saxpy_vec4(int n4, float a, const float4* __restrict__ x, float4* __restrict__ y) {
    const int stride = gridDim.x * blockDim.x;
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n4; i += stride) {
        float4 xv = x[i];
        float4 yv = y[i];
        yv.x = fmaf(a, xv.x, yv.x);
        yv.y = fmaf(a, xv.y, yv.y);
        yv.z = fmaf(a, xv.z, yv.z);
        yv.w = fmaf(a, xv.w, yv.w);
        y[i] = yv;
    }
}

static void check(cudaError_t e, const char* what) {
    if (e != cudaSuccess) { printf("CUDA error at %s: %s\n", what, cudaGetErrorString(e)); exit(1); }
}

int main(int argc, char** argv) {
    const int  n   = (argc > 1) ? atoi(argv[1]) : (1 << 24);
    const int  rep = (argc > 2) ? atoi(argv[2]) : 20;
    const float a  = 2.5f;
    const size_t bytes = (size_t)n * sizeof(float);

    float *x, *y;
    check(cudaMalloc(&x, bytes), "malloc x");
    check(cudaMalloc(&y, bytes), "malloc y");
    check(cudaMemset(y, 0, bytes), "memset y");
    check(cudaMemset(x, 1, bytes), "memset x");

    const int threads = 256;
    const int blocks  = (n + threads - 1) / threads;

    cudaEvent_t e0, e1;
    cudaEventCreate(&e0); cudaEventCreate(&e1);

    auto bench = [&](const char* name, void (*kern)(int, float, const float4*, float4*), bool vec,
                     int grid) {
        for (int r = 0; r < rep; ++r) {                       // 每轮重新计时，含 warmup 轮
            cudaEventRecord(e0);
            if (vec) kern(n / 4, a, (const float4*)x, (float4*)y);
            else     saxpy_scalar(n, a, x, y);
            cudaEventRecord(e1);
            cudaEventSynchronize(e1);
            float ms = 0.f; cudaEventElapsedTime(&ms, e0, e1);
            if (r == rep - 1)
                printf("%-8s grid=%d  time = %8.3f ms   traffic = %6.1f MB   %7.1f GB/s\n",
                       name, grid, ms, 3.0 * bytes / 1e6, 3.0 * bytes / (ms * 1e-3) / 1e9);
        }
    };

    // 版本 A：一次性大 grid（一线程一元素）
    bench("scalar", nullptr, false, blocks);
    // 版本 B：grid-stride + float4，grid 选成"足够覆盖设备"的固定值
    bench("vec4",   saxpy_vec4, true, 1024);

    check(cudaGetLastError(), "launch");
    cudaFree(x); cudaFree(y);
    return 0;
}

• 【代码做什么？】 两个 kernel 计算同一个 y = a·x + y（SAXPY，每个元素 1 次 FMA = 2 FLOP）：版本 A 每个 CUDA 线程处理 1 个元素、每线程每流只搬 4 B；版本 B 用 grid-stride 循环（固定 grid，让每个线程处理多个元素，便于覆盖不同设备规模）+ float4 向量化访存（每线程每流 16 B，正好凑满 128 B 的一次内存事务/扇区），并用 __restrict__ 告诉编译器 x/y 不重叠以允许更激进的访存调度。计时用 CUDA event（cudaEventElapsedTime），它测的是设备上 kernel 的实际执行时间，不受主机侧启动开销与异步性影响。

• 【并行机制与性能解说】

  • 硬件执行模型：threads=256 定义 block（线程块）；每个 block 被整体调度到一个 SM（Streaming Multiprocessor）上，block 内的线程以 warp（32 个线程） 为调度单位执行 SIMT。版本 A 的 if (i < n) 只对最后不满一个 warp 的边界产生分支发散，主体部分无发散。版本 B 的 float4 让每个线程在一条指令里处理 4 个元素，指令数降到 1/4，同时把每线程每流的字节数从 4 B 提高到 16 B ⇒ 更少的 warp 就能喂满内存系统（提高 MLP / memory-level parallelism）。
  • Work / Span / 并行度：N = 2²⁴ = 1.68×10⁷ 个元素。
    • Work W = N 次 FMA（= 2N FLOP = 33.6 MFLOP）。
    • Span：把执行抽象成”一次访存延迟 + 一次 FMA”的链，整个 kernel 的 Span 取决于同时驻留的线程数 T_res：Span ≈ Θ(N / T_res)。取 80 个 SM × 每 SM 2048 线程 = T_res = 163,840，则 Span ≈ 102 个”批次”。
    • 并行度 = W/Span ≈ T_res = 1.6×10⁵。结论极其重要：并行度对这块 GPU 而言绰绰有余（比可驻留线程数高两个数量级余量），所以瓶颈绝不在”并行度不够”，而在内存系统。这正是”先算 work-span 再动手”的价值：它把”要不要加并行度”这个问题直接排除了，使你把注意力放到 float4 / 访存事务 / 占用率上。
  • 瓶颈分析（流量下界）：SAXPY 读 x、读 y、写 y ⇒ 3N×4 B = 192 MiB = 201 MB。若该 GPU 的可持续带宽为 1.4 TB/s，则时间下界 = 201 MB / 1.4 TB/s ≈ 144 µs。若实测版本 A 只跑出 ~45% DRAM 利用率（≈ 0.63 TB/s ⇒ ~320 µs），而版本 B 达到 ~90%（≈ 1.26 TB/s ⇒ ~160 µs），则加速比约 2×，且完全由”每个线程请求的字节数/事务效率”解释，与算力无关（算力时间 33.6 MFLOP / 19.5 TFLOP/s ≈ 1.7 µs，比内存时间小两个数量级）。算术强度 = 2 FLOP / 12 B = 0.167 FLOP/B，在该 GPU 的拐点（19.5e12 / 1.4e12 ≈ 14 FLOP/B）左侧约 84 倍，是彻头彻尾的带宽受限内核。
  • profiling 反馈循环（本示例的”元”信息，直接来自讲义 slide 45–46 的范式）：讲义的循环是——运行并剖析 → 得到统计量（Timing / SM util / DRAM util / L2 hit rate）→ 让优化者（人或 LLM）”反思什么在拖慢程序” → 据此做一处修改 → 再跑再测。把它套到本示例上就是一条可复现的诊断链：
    1. SM util 42% + DRAM util 89% ⇒ 不是算力不够，是访存已经打满——在这种情况下任何提高 ILP、增加算力的修改都不会有效，唯一有效的方向是减少字节数或提高事务效率（float4、__ldg、融合、降低精度）。
    2. L2 cache hit rate 68% 提示还有约 1/3 的 L2 请求下行到 DRAM；如果这是重复访问造成的（如分块/多趟算法里的中间结果），则融合（fusion）或分块（tiling）是正确手段（对应 §3.1 的 chunked 版本）。
    3. 若 SM util 高而 DRAM util 低 ⇒ 反过来，计算/指令受限（分支发散、特殊函数单元、共享内存 bank 冲突），此时”减少访存”只会浪费精力。
  • 与自动搜索/自动调优的接口：讲义指出，把优化目标写成”测量量“之后，优化本身就可以被搜索：Halide 自动调度器把”调度空间”枚举出来，用一个小型 MLP 代价模型（输出 27 个系数供手写代价模型使用）在 166 秒内评估 140 万个 schedule，其质量已”compares to best known human schedules”；而人类专家要达到同等水平通常要花 10–50 分钟（讲义 slide 43 的时间—吞吐曲线）。这说明：可信的、快速的性能测量是自动优化的前置条件——没有它，搜索就没有目标函数。

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

本节的机器模型（贯穿全节）：16 核 × 3.0 GHz × 8 宽 SIMD × 2（FMA）= 768 GFLOP/s（fp32）；可持续 DRAM 带宽 20 GB/s；DRAM 延迟 ~100 ns；L1 64 B/cycle；64 B cache line。 拐点 P_peak/BW = 768/20 = 38.4 FLOP/Byte。

4.1 算例 A：算术强度与 Roofline 的完整数值推理（3×3 box blur）

假设参数：W = H = 1024（1,048,576 像素）；float 4 B。

结论与可执行判据：

1. 三个版本全都落在 Roofline 斜段（AI 远小于 38.4），所以唯一的有效优化方向是”减流量”：把两趟变成分块融合，流量从 16.80 MB 降到 8.40 MB ⇒ 理论加速比 2.0×，这与讲义”两趟实现算术强度低 2 倍（2x lower arithmetic intensity）”的结论完全一致。
2. 二维（不可分离）版本的 AI 是两趟朴素版的 3 倍（2.25 vs 0.75）——虽然它做了更多算术（9 MAC vs 6 MAC），但因为不产生中间流量，在带宽受限的机器上反而可能更快。这是”更多计算、更少搬运可以更快”的经典反直觉结论，也正是 Roofline 要教的判断方式（讲义原文即谓两趟是”loads/stores to tmp_buf are overhead — this memory traffic is an artifact of the two-pass implementation: it is not intrinsic to computation“）。
3. “固有流量”（intrinsic）的界定：讲义明确指出，blur 的固有带宽需求就是”读一遍输入、写一遍输出”。任何超出这一点的流量都是实现的人工产物（artifact），是优化的首要目标。分块/融合（CHUNK 增大）让实际流量单调趋向固有流量（(CHUNK+2)/CHUNK → 1）。
4. 验证方式：算出的 20 GB/s 是否可信？用程序打印的 GB/s 与 perf stat 的 IMC/uncore 字节数交叉验证；用 STREAM 类内核实测本机可持续带宽作为分母。若实测 GB/s 已达可持续带宽的 90%+，则本分析闭合；若只有 30%，说明还有延迟/MLP/同步问题没被这个带宽模型解释，需要回到 §2.5 的高水位实验。

4.2 算例 B：T = 2N²/P + P ——并行度、屏障与”加核到什么时候停止”

来自 CS149 thoughtprocess 讲义的两阶段图像处理例子：步骤 1（亮度 ×2，可完全并行）、步骤 2（求全图平均，需要部分和 + 合并）。串行时间 ≈ 2N²；并行实现的时间是”每个处理器 N²/P（步骤 1）+ 每个处理器 N²/P（步骤 2 的部分和）+ P（串行合并部分和）”：

T(P) = 2N²/P + P
S(P) = 2N² / T(P) = P / (1 + P²/(2N²))
E(P) = S(P)/P = 1 / (1 + P²/(2N²))

取 N = 1024（N² = 1,048,576）：

解读：S(P) 在 P = √2·N ≈ 1448 处取极大值 724，之后随 P 增加而下降。这条曲线同时包含了两种效应的竞争：分子上的 P 来自”可并行部分随 P 缩短”，分母里的 P² 来自”串行合并/屏障的绝对成本随 P 线性增长“（P 项）。在 16 核机器上（P=16），效率 99.99%，并行几乎免费；但把同一份代码搬到 2048 核上，效率掉到 33%，加速比回落到与 1024 核相同的 682.7×、并低于峰值 724×——即”加了一倍的核，一点没变快”。 这与讲义 slide 35 的原话一致：“speedup → P when N » P”，以及”并行算法的开销：合并部分和”。它是 Amdahl 定律的一个具体化——只不过这里的”串行部分”随 P 增长（屏障/归并/同步），而 Amdahl 的经典形式把它当作常数。

与 Amdahl 的对照算例（f = 恒定串行比例）：

• f = 1%，P = 64：S = 1/(0.01 + 0.99/64) = 1/0.025469 = 39.26，效率 61.3%；用 Karp-Flatt 度量反推串行比例：e = (1/S − 1/P)/(1 − 1/P) = (0.025469 − 0.015625)/0.984375 = 1.00% ✔（自洽）。
• Summit 超算的极端算例（讲义 slide 37）：27,648 块 GPU × 5,376 ALU/GPU = 148,635,648 个 ALU，即机器可同时执行 1.486 亿次单精度运算。若应用有 0.1% 串行，则 S_max = 1/0.001 = 1000——这台 1.5 亿 ALU 的机器最多只能带来 1000 倍加速。把它换算成”机器白买了多少”：T(P) ≈ 0.001·T(1)，这段时间内可提供的 ALU-时间容量为 P·T(P) = 1.486e8 × 0.001·T(1) ≈ 1.486e5·T(1)，而真正有用的工作量只有 T(1)（其中 0.999·T(1) 被摊在 P 个 ALU 上、0.001·T(1) 只由 1 个 ALU 完成），因此按时间平均的算力利用率仅约 1/1.486e5 ≈ 6.7×10⁻⁶（0.0007%），99.9993% 的 ALU 时间容量被浪费。这就是”小的串行段会限制大规模并行机器上的加速比“的定量含义：瓶颈不是机器不够大，而是串行段让机器无事可做。
• Gustafson 对照：同一 f = 0.01、P = 64，若问题规模随机器放大（弱扩展），S_scaled = P − f(P−1) = 64 − 0.63 = 63.37，效率 99.0%。同一个 f，强扩展只有 61% 效率、弱扩展有 99% 效率——这解释了为什么 HPC 社区的实测加速比通常远好于 Amdahl 强扩展预测。

4.3 算例 C：带宽、延迟与 Little’s Law ——”我需要多少并发才能打满带宽”

公式：需要的在飞字节数（bytes in flight）= 带宽 × 延迟。代入 BW = 20 GB/s、L = 100 ns：

20e9 B/s × 100e-9 s = 2000 B  ≈ 31 条 64 B cache line  ≈ 16 条 128 B line

含义：要让 DRAM 带宽饱和，内存系统里必须同时有约 31 条 cache line 的请求在飞。若每个线程只有 1 个未完成的缺失请求（MLP = 1），你至少需要 31 个活跃线程；如果每线程 MLP = 4（例如 float4 或循环展开后 4 个独立缺失），需要 约 8 个线程。这解释了 §3.3 中 GPU 版本 B 的行为：float4 把每线程的字节数从 4 B 提到 16 B，等价于把”每事务有效载荷”提高 4 倍，从而用更少的 warp 就能达到同样的在飞字节数。

与”延迟 vs 带宽”的判据对应：

• 实测带宽 ≪ 机器可持续带宽，且增加并发（线程/SIMD 宽度/展开）能线性提速 ⇒ 你受限于延迟/MLP，解决办法是更多并发（预取、展开、软件流水）。
• 实测带宽 ≈ 机器可持续带宽，增加并发不再提速 ⇒ 你受限于带宽，唯一出路是减少流量（分块、融合、降低精度、压缩）。
• 这两者的区分是整个性能分析中最重要的一次分岔，而它必须用测量（计数器字节数 + 并发扫描）而不是直觉来判定。

4.4 算例 D：测量本身的不确定度 ——”你的 3% 提升是真的吗”

设某内核单次运行时间服从均值 μ = 3.20 s、标准差 σ = 0.40 s（σ/μ = 12.5%，典型的多线程抖动水平）。均值估计的标准误差为 SE = σ/√n：

结论：

1. 要检测 1% 的性能差异，需要约 156 次重复运行（n = (σ/(目标SE))² = (12.5/1)² ≈ 156）。跑一次就宣布”我优化了 3%”是不成立的——单次测量的 1σ 不确定度就是 12.5%。
2. 应当报告中位数 + min + max（或分位数），而不是均值：多线程计时分布右偏（偶发的调度/争用事件拖长尾巴），均值会被离群值拉偏，而 min 最接近”无干扰性能”。这正对应 §3.1 代码打印三个统计量的原因。
3. 采样剖析的误差有同样的性质：§2.3 的表格给出 ±1 个百分点 → 9604 个样本 → 1 kHz 下约 9.6 s。“这个热点占 X%” 必须附上样本数，否则无法判断结论是否显著。
4. 测量误差的系统性来源必须逐条排除：计时器分辨率与开销、被优化掉的死代码（DCE）、CPU 频率/睿频波动（用 cycles/time 反算实际频率）、首次触碰的页错误与 NUMA first-touch、冷 cache（是否 warmup）、超线程共享执行单元导致的干扰、以及同机其它进程的噪声邻居。

4.5 模型选择速查表：三个上界各管什么

4.6 算例 E：把 profiler 数字翻译成结论（综合演练）

给定一份 profiler 报告（形式与讲义 aiperfoptimization slide 46 的反馈循环一致；此处为用于演示诊断流程的示例数字，并非讲义原文数据）：

Timing: 32 ms          SM util: 42%          DRAM util: 89%          L2 hit rate: 68%

推理链：

1. DRAM util 89% 是第一优先信息：访存子系统几乎打满，说明程序已经用掉了近 9 成的可用带宽。任何”提高算力利用率 / 增大 ILP / 换更快指令”的修改都不可能带来大收益——因为瓶颈不在计算侧。
2. SM util 42% 与上一条互证：SM 利用率不高不是”算力浪费”，而是在执行单元等待数据（访存受限时的典型形态）。不要去追这 42%。
3. L2 hit rate 68% 指出改进方向：有约 1/3 的 L2 请求下行到 DRAM。若这些下行是重复访问造成的（多趟/多内核之间的中间结果、未复用的邻域），则融合或分块能直接减少下行量（正是 §3.1 chunked 的手法，收益约 2×）；若是一次性流式访问（每个字节只读一次），则下行无法避免，改进只能来自降低字节需求本身（更窄的数据类型、只读必要字段、压缩）。
4. 给出上界估算：若把 DRAM 流量减半而其它不变，时间从 32 ms 降到约 32×(0.11 + 0.89/2) = 17.8 ms（把 11% 视为与 DRAM 无关的、不可压缩的部分）。这个 1.8× 就是”减流量”路线的 ceiling；如果业务要求 5×，那么必须在算法层面减少工作或减少精度，而不是继续在 kernel 里微调。
5. 反过来验证：跑一个高水位实验——”把大部分计算删掉但保留全部访存”。若时间几乎不变（例如 32 ms → 29 ms），则访存主导成立；若掉到 12 ms，则之前对 SM util 42% 的解读就有问题，需要重新用 stall 分解（stalled-cycles-frontend/backend）检查前端/后端停顿的来源。模型给出假设，激发实验负责证伪。

5. 关键要点

1. 测量优先，且测量必须可信。 讲义的原话是”永远、永远、永远先写最简单的并行版本，然后测量性能，看看你站在哪里“。可信的测量意味着：用墙钟（不是 clock()）、做 warmup、重复多次并报告中位数与 min/max（检测 1% 差异需约 156 次运行）、防止死代码消除、明确基线（与最优串行实现比，而不是与并行算法的单核运行比）。测不准的优化等于随机游走。

2. 性能问题只有四类来源，先分类再优化：计算、带宽/延迟、同步、并行度。 用三个工具分诊：work-span（W/S 够不够 → 加核有用吗）、算术强度 / Roofline（AI vs P_peak/BW → 该减流量还是减计算）、PMU 计数器 + 高水位实验（IPC、cache 命中率、DRAM 字节数、stall 分解 → 延迟界还是带宽界、是不是同步主导）。分类错了，优化就是白工（例：在 DRAM util 89% 的内核上调 ILP）。

3. 算术强度是并行性能的单一最重要定量指标，Roofline 是它的可视化。 可达性能 = min(P_peak, AI × BW)，拐点 = P_peak/BW。本笔记的 CPU 机器模型下拐点为 38.4 FLOP/B，而 blur 的 AI 只有 0.75–2.25 ⇒ 只能拿到峰值的 1.95%–5.9%；SAXPY 更低（AI = 0.167）——即便换到 GPU 机器模型（19.5 TFLOP/s ÷ 1.4 TB/s ⇒ 拐点约 14 FLOP/B），它同样低两个数量级。跨越拐点的唯一途径是数据复用（分块、融合、缓存阻塞）或降低每字节的算术需求；在拐点左侧做的所有算力优化都收益接近于零。

4. “实现”（implementation）而不是”抽象”（abstraction）决定性能，所以要会掀开后厨。 ISPC 的同一个 foreach 抽象，交错分配编译成一条 packed load（vmovaps）、分块分配退化成昂贵的 gather（vgatherdps）；同一份 A[i,j] 更新代码，把锁从”每元素一次”改成”每线程每迭代一次”就差一个数量级（讲义 slide 68→69）。性能分析必须一直问到”硬件上实际发生了什么”这一层（看汇编、看计数器、看访存模式），否则你优化的是心里的模型，不是机器上的程序。

5. 测量与建模的终极目的是”自动化”——而自动化质量的上限由测量的可信度决定。 把性能写成可测量的目标函数之后，优化就可以被搜索：Halide 自动调度器用小型 MLP 代价模型在 166 秒内评估 140 万个调度，质量已与最佳人工调度相当（人工通常要 10–50 分钟一个）；基于 profiling 统计量的”反思—修改—再测量”循环（包括 LLM agent）也是同一范式。但所有这些方法的输入都是 profiler 的四个数字——测量不可信，搜索就会朝错误方向收敛。“建立高水位、分清瓶颈”这项人类技能，在自动优化时代不是被取代，而是变成自动优化的地基。

6. 常见陷阱与注意事项

• 测错了”时间”。 用 clock() / getrusage 的 CPU 时间与墙钟时间算加速比，在多线程下会得出错误甚至负的加速比（8 线程时 CPU 时间约为墙钟的 8 倍）。并行加速比只能用墙钟时间。 另外要留意：omp_get_wtime 不保证全局同步时钟，跨节点计时要用 MPI_Wtime。
• 单次测量下结论。 多线程程序的计时分布右偏、抖动可达 10% 以上；单次运行的 1σ 不确定度可能就是 12.5%（§4.4）。必须 warmup + 重复 + 报告中位数与 min/max；检测 1% 的差异需要约 156 次运行。同理，采样剖析的百分比必须附样本数：0.3% 的热点在 1000 个样本里只有 3 个样本，泊松相对误差约 58%，这个结论不可用。
• 被测代码被优化掉（DCE）。 -O3 下未被使用的计算结果会被整段删除，于是你”测”出一个 0 ms 的内核。必须消费结果（打印校验和、写 volatile、内联汇编屏障）。反过来的错误同样常见：为了防 DCE 而加 volatile 到热循环内的每一次数组访问上，会把向量化和重排全部禁掉，测出的不是原程序的性能。
• 基线选错、问题规模选错。 讲义明确点名的两个陷阱：①“把并行程序的加速比对比于并行算法在单核上运行”（原文：easier to make yourself look good）——并行算法往往做更多总工作（如红黑着色比 Gauss-Seidel 收敛更慢），这样做会凭空造出加速比；②固定问题规模——32 处理器上跑 258×258 网格得到 “No benefit! (slight slowdown)”，而 1K×1K 网格就正常扩展，因为小问题对机器来说太小、通信计算比太高。报告加速比时必须写清基线、问题规模、线程数、以及是强扩展还是弱扩展。
• 把伪共享/同步问题误判成带宽问题（或反之）。 两者症状相似（加线程不变快甚至变慢），但证据完全不同：伪共享/一致性争用的特征是 DRAM 字节数低而一致性事务/line 迁移多、时间随线程数恶化；真实带宽瓶颈的特征是从内存控制器读到的字节数已逼近机器上限。用 perf stat -e LLC-load-misses,... 与 perf c2c report（直接指出争用的地址）区分，而不是凭”感觉像内存问题”。另外记住修伪共享的正确方向是把”共享写”改成”私有写 + 归并”（reduction），padding 只是次优解。
• 忽略”测量者效应”与硬件状态。 插桩剖析会改变被测程序（mcount 让小函数时间虚高、不支持共享库语义）；PMU 事件数超过物理计数器时会多路复用外推（perf stat 的 [xx%] 列必须看）；频率/睿频、SMT 共享执行单元、NUMA first-touch、冷 cache、页错误、噪声邻居都会污染结果。跨机器比较原始计数没有意义，只能比较比值（IPC、命中率、字节/秒、% of peak）。
• 只相信一个模型，不做激发实验。 Roofline 上界乐观、Amdahl 的 f 恒定是理想化、work-span 假设零通信。讲义的原话是”计算、访存与同步几乎从不完美重叠，因此总体性能很少完全由其中之一决定”。正确姿势是：模型提出假设 → 高水位实验（加算 / 全改 A[0] / 去锁 / 去算保访存）负责证伪 → 计数器给出机制级证据。三者不一致时，以实验为准，并去解释模型为何失效。

7. 思考题（带答案）

问题 1（测量方法论）：某同学在 16 核机器上实现了一个 5 点模板计算，报告”8 线程比 1 线程快 13.8×，因此我的并行实现是超线性的”。他的实验方法是：单线程版本用 clock() 计时（因为”CPU 时间更准”），多线程版本用 omp_get_wtime() 计时；每个配置只跑一次；单线程版本不 warmup，多线程版本 warmup 一次。请指出这份报告中的全部方法论问题，说明每个问题会朝哪个方向污染结论，并给出正确的实验设计。

【答案】 问题与污染方向：

1. 两个版本用了不同的计时域（clock() 是 CPU 时间，omp_get_wtime() 是墙钟）。clock() 在单线程下与墙钟接近（差别是系统调用/I/O），所以单线程基线”看起来”问题不大；但一旦代码在单线程下也用了任何库内线程（如 BLAS 或 OpenMP runtime 的 helper），clock() 会高估基线时间，从而虚增加速比。更严重的是，两个域不可比 ⇒ 整个加速比在量纲上就不成立。方向：可能虚增。 正确做法：两个版本都用同一个墙钟计时器（omp_get_wtime / steady_clock）。
2. 只跑一次。多线程计时分布右偏、抖动可达 10%+（§4.4）。单次测量的 1σ 不确定度约 12.5%。要检测 1% 的差异需约 156 次重复。方向：结论完全不可信，可能虚增也可能虚减，量级上足以解释 13.8× 与 8× 之间的全部差距。 正确做法：warmup + 至少数十次重复 + 报告中位数与 min/max。
3. warmup 不对称。单线程不 warmup 会包含首次分配页错误、first-touch NUMA 放置、冷 cache、首次 JIT/动态库加载等一次性成本，从而高估单线程时间、虚增加速比。多线程 warmup 一次也不够（第一次并行区会包含线程创建与栈分配）。方向：虚增。 正确做法：两个版本同样 warmup，并把 warmup 轮从统计中剔除。
4. 超过核数的”加速比”应当首先被怀疑是硬件状态差异，而不是算法优越。16 核上正确实现的 5 点模板的强扩展上限就是 16（且实际会因带宽受限远低于 16）。13.8× 本身没超过 16，所以不是数学上不可能，但在内存带宽受限的模板计算上达到 86% 效率需要额外解释。必须提供的证据：①实测 GB/s 是否逼近机器可持续带宽（若是，说明 8 线程已打满带宽，1 线程只跑出 1/8 带宽是完全可能的，此时 8× 就是上限而 13.8× 需要解释）；②perf stat 的 IPC 与 LLC 命中率在两种配置下是否发生了质变（例如单线程工作集超出 LLC 而多线程各自分块后命中 LLC ⇒ 这才是真正的”超线性”，且属于工作集变小的效应，而不是并行本身带来的）。
5. 缺少基线说明：加速比的基线必须是”最优串行实现“（讲义原文的 pitfall），且要说明是强扩展还是弱扩展。 正确的实验设计：(a) 统一墙钟计时器；(b) 两版本对称 warmup；(c) 每个配置 ≥50–200 次重复，报告中位数/min/max/标准差；(d) 线程数扫描 1,2,4,8,16 并画出实测加速比 vs Amdahl 上界及实测 GB/s两条曲线；(e) 用 perf stat 记录两配置的 cycles,instructions,LLC-load-misses 与实测频率；(f) 用 STREAM 类内核测出本机可持续带宽作为分母；(g) 报告基线、问题规模、是否强扩展。

问题 2（Roofline 应用与决策）：机器参数：16 核 × 3.0 GHz × 8 宽 SIMD × 2（FMA）= 768 GFLOP/s，可持续 DRAM 带宽 20 GB/s。你的内核每次处理一个 1024×1024 的 float 图像（4 MB），当前实现是”两趟”：第一趟读输入、写中间缓冲（4 MB）；第二趟读中间缓冲、写输出（4 MB）。请回答：(a) 算术强度是多少（内核每像素做 6 次乘加）？(b) 带宽下界时间是多少？(c) 若 profiler 显示当前耗时 3.4 ms，你的第一嫌疑是什么？(d) 把两趟融合成分块实现后流量降到 8.4 MB，预期时间与加速比是多少？(e) 若融合后实测只到 1.1 ms 而非预期值，你会用哪些测量手段继续定位？

【答案】 (a) 流量：4.21 MB（读输入，含 (W+2)×(H+2) 的边界）+ 4.20 MB（写中间）+ 4.20 MB（读中间）+ 4.19 MB（写输出）= 16.80 MB。FLOP：6 次乘加/像素 × 2 = 12 FLOP/像素 × 1.048576e6 像素 = 12.58 MFLOP。AI = 12.58/16.80 ≈ 0.75 FLOP/Byte（远低于拐点 38.4 ⇒ 带宽受限，理论上限 0.75 × 20 = 15.0 GFLOP/s，仅峰值的 1.95%）。 (b) 带宽下界 = 16.80 MB / 20 GB/s = 0.84 ms。算力时间 = 12.58 MFLOP / 768 GFLOP/s = 16.4 µs，比内存时间小 51 倍 ⇒ 纯带宽/算力模型下这是内存受限内核。 (c) 实测 3.4 ms 是带宽下界的 4.05 倍 ⇒ 第一嫌疑不是“算力不够”（那只会让它更快），而是实际可达带宽远低于 20 GB/s，原因按概率排序：①访存模式/MLP 不足（每线程未完成缺失太少，受延迟而非带宽限制；见 §4.3：打满 20 GB/s 需要约 2000 B 即 ~31 条 cache line 在飞）；②中间缓冲的实际 DRAM 往返比模型更多（写分配导致”写”也要先读一次：若 tmp 写是 write-allocate 且未用非临时写，则额外多出 4.20 MB 读，总流量升到 21 MB ⇒ 下界 1.05 ms，仍不足以解释 3.4 ms）；③同步/屏障或负载不均（parallel for 每趟一次屏障，若线程数少或行数分配不均，这部分不可压缩）；④多线程未生效或超线程争用（若有效线程数实为 4，则带宽下界 3.36 ms 恰好与 3.4 ms 吻合——这是一个必须首先排除的假设）。 (d) 融合后流量 8.40 MB ⇒ 下界 0.42 ms；相对当前实现的理论上限加速比 = 16.80/8.40 = 2.0×，即理论时间从 0.84 ms 降到 0.42 ms。若当前实现已达其自身的带宽下界（0.84 ms），融合会把时间也降到 0.42 ms（2×）；但你实测是 3.4 ms，所以融合本身不会自动带来 8×——它只把”流量”这一项减半，而当前的主要损失在流量之外。 (e) 融合后实测 1.1 ms（比下界 0.42 ms 差 2.6 倍）时的定位手段：

1. 先确认有效并行度：打印 omp_get_num_threads() 与 OMP_NUM_THREADS，并用 perf stat -e task-clock / /proc/<pid>/status 核对实际使用的 CPU 数；排除”线程没起来”或”绑核到同一物理核”。
2. 量出真实带宽：用 perf stat -e LLC-load-misses 或 IMC/uncore 事件读从内存控制器实际搬运的字节数，与时间相除得到实测 GB/s，与 STREAM 类内核测得的机器可持续带宽比较。若实测已达 90%+，则程序其实已经到位，20 GB/s 的假设偏乐观（真实机器可能只有 8–12 GB/s）——是模型需要修正，而不是程序。
3. 区分延迟界与带宽界：做并发扫描——固定问题、扫线程数 1→16，同时扫循环展开/float4 向量化。若带宽随并发线性上升到某点后走平，走平点即带宽界；若始终线性上升，说明仍是延迟/MLP 界，应加预取或展开。
4. 高水位实验定位同步成本：删掉屏障/合并（结果允许暂时错误）测时间；若明显变快，说明同步/负载不均占了大头。同时用 schedule(static) vs dynamic 对比，检查负载不均。
5. 检查 cache 与页行为：perf stat -e L1-dcache-load-misses,LLC-load-misses 看命中率；确认分块尺寸真的让中间缓冲常驻（CHUNK=16 时 18×1024×4 = 73.7 KB 是否小于该机器的 L2 私有容量）；用 numactl --membind / first-touch 确认没有远地 NUMA 访问（远地带宽可能只有本地的一半）。
6. 最后才怀疑指令层：看 stalled-cycles-frontend/backend、IPC，以及反汇编确认内层循环确实被向量化（没有向量化则单核算力降 8 倍，但在此内核里算力远非瓶颈，所以这一项优先级最低——这正是”先分类再优化”的价值）。

问题 3（work-span、伪共享与”加线程为什么没用”）：下面的代码在 4 核机器上跑 T = 4 个线程、每线程 10⁸ 次自增，实测 11.2 秒；把数组元素改成 alignas(64) 各自独占一行后实测 0.28 秒。(1) 用 work-span 的语言解释为什么第一版的”并行度”实际不到 1；(2) 说明 11.2 s 与 0.28 s 分别与什么硬件量对应，并给出量级估算；(3) 是否应该继续加线程到 8？(4) 如果这段自增是”统计各类别事件的计数器”（每处理一条记录就给对应类别的计数器加 1），且类别数 = 1000，请给出一个比 padding 更好的修改方案，并说明它为什么更好。

// counter_frag.c —— 为聚焦问题而缩到最小的片段（不含计时与线程创建）
// 编译: gcc -O3 -pthread counter_frag.c -o counter_frag
#include <pthread.h>
#include <stddef.h>
#define NTHREADS 4

long long cnt[NTHREADS];   // 全局数组，元素相邻（全部落在同一条 64 B cache line 内）
void* worker(void* arg) {
    int t = *(int*)arg;
    for (long long i = 0; i < 100000000LL; ++i) cnt[t]++;   // 自增
    return NULL;
}

【答案】 (1) Work-Span 分析：Work W = T × 10⁸ = 4×10⁸ 次自增（每次是 1 次读-改-写）。问题在于 Span：4 个 cnt[t] 元素位于同一条 64 B cache line 内，x86 的一致性协议以 cache line 为所有权粒度，因此线程 t 写 cnt[t] 必须独占整条 line；任一时刻只有一个 core 能持有该 line 的 M 态，自增被强制串行化。串行化的单位成本是一次 line 迁移（cache-to-cache transfer + 失效 + 所有权获取），实测约 数十到约 100 ns。于是 S ≈ W × L_transfer = 4×10⁸ × （数十 ns） ⇒ 用 11.2 s 反推 L_transfer = 11.2 s / 4e8 = 28 ns，即每次自增平均 28 ns——完全符合”每次自增 = 一次跨核 line 迁移”的预期。因此并行度 = W/S = 1/L_transfer ≈ 3.6×10⁷ 次/秒，与线程数 T 无关；等价说法是：真实并行度 ≈ 1（4 个线程在队列里轮流拿同一条 line）。加线程只会让 line 搬家更频繁，不会提高吞吐。

(2) 11.2 s 对应”一致性/延迟串行化”，其控制量是 line 迁移延迟（~28 ns 量级），而不是带宽或算力：DRAM 字节数极低（数据本来就在 cache 里），ALU 利用率极低（每次自增之间要等几十 ns）。0.28 s 对应”ALU/流水线极限（含私有 line 的 L1 命中）”：4×10⁸ 次自增 / 0.28 s = 1.43×10⁹ 次/秒；分摊到 4 个核、每核 3.0 GHz，即每核约 0.12 次自增/周期。这个数偏低（一个 inc 序列在 L1 命中的情况下本可接近 1 次/周期），原因是 cnt[t]++ 是”load–add–store”且相邻迭代存在循环携带依赖（同一个计数器），每轮至少要等一次 L1 命中延迟（~4–5 周期）+ 一次 store-to-load forwarding，再加上 volatile/编译屏障（若原代码用了）阻止了流水化。量级结论：11.2/0.28 = 40×，与”跨核 line 迁移 ~28 ns 对比 L1 命中约 1–2 ns”的 ~20–30× 比值同量级，多出来的部分来自协议开销（请求转发、目录/侦听、失效确认）。

(3) 不应该。 因为 W/S 已经与 T 无关（≈3.6×10⁷ ops/s，由 line 迁移延迟决定），加线程到 8 的预期结果是”不变或更慢”——所有 8 个线程仍在争同一条 line，只会增加协议流量与排队。这与 §3.2 的分析完全一致：伪共享下”并行度”从 T 掉到 1 以下，加核无用。正确的下一步不是加线程，而是消除争用（padding、或按 (2)(4) 的方案改造算法）。

(4) 更好的方案：按线程/块做”私有直方图 + 最后归并”（privatization），而不是给 1000 个计数器逐个填充到 64 B。

• 做法：每个线程先在自己的私有局部直方图（long long local[NBINS]）上累加，全部处理完后，再在临界区里做一次 atomic/reduction 把 T 份局部直方图按 bin 相加（归并成本 O(NBINS × T)，与 10⁸ 次自增相比可忽略）。
• 为什么比 padding 更好：①padding 的容量代价是 64 B × 1000 = 64 KB（每线程！4 线程 256 KB），会溢出 L1（32–48 KB）、挤占 L2，并把直方图从”L1 常驻”变成”L2/L3 往返”，反而增加真实带宽需求；私有直方图只需 8 B × 1000 = 8 KB/线程，稳稳装进 L1。②padding 只消除了跨核争用，但每次自增仍是”共享可写位置”的读-改-写，且若同一线程内多个 bin 被高频更新仍有写端口/转发压力；私有化则让每次自增都命中本核 L1 中自己独占的行，并把跨核通信压缩成一次批量归并。③这是本课程反复出现的通用范式：“共享写”→”私有写 + 归并”，与讲义把”每个 (i,j) 加锁”改成”每线程累加 myDiff、每次迭代只加锁一次”（thoughtprocess slide 68→69）是同一个手法的两种表现。
• 正确性注意：计数器的最终值必须通过一次原子/同步的归并得到（#pragma omp atomic 或 #pragma omp critical，或 #pragma omp parallel for reduction），归并点是唯一的同步需求；同时要保证归并发生在线程全部完成本地累加之后（隐式 barrier）。