Lecture 24: AI in System Design

目录 · ← l22 · l24 →

Lecture 24: AI in System Design

1. 章节标题与概述

Lecture 24: AI in System Design（用机器学习替换系统里的手写启发式：learned index、learned cost model、学习型调度与自动调优）

• 本讲核心问题：过去四十年，计算机系统的每一个”该怎么做”的决定——缓存该替换哪一行、DRAM 控制器该先服务哪个请求、编译器该给循环选多大的分块、运行时该把迭代交给哪个 lane / 哪个线程、推理框架该把多少个算子融合成一个 kernel——都是人写死的启发式（heuristic）。本讲问的是两件事：（1）当数据（遥测、访存轨迹、性能计数器、编译日志）变得廉价而模型变得强大时，能不能把这些启发式换成一个从数据里学出来的模型（learned policy / learned cost model）？ （2）要让这种”用 AI 设计系统”的做法真正落地，代价是什么？——最关键的代价是：学习到的策略必须在它所优化的那台系统上、以远低于它想省下的开销运行（例如手机 SoC 上一次访存的能耗预算只有约 10 nJ），于是学习型策略的推理成本（推理 FLOP、模型体积、SRAM 占用、决策频率）本身成了一个必须做 work-span / Roofline / 能耗分析的系统设计问题。

• 涉及的主要硬件/软件机制：
  • 硬件侧：HBM（High Bandwidth Memory）与 3D 堆叠 + TSV（把内存搬到处理器旁边，1024-bit/stack 接口，H100 达 3.2 TB/s）；DRAM bank / row buffer / burst mode 与内存控制器的请求调度（FR-FCFS 启发式、每 bank 请求队列、请求合并）；脉动阵列（systolic array）（TPU 风格稠密矩阵乘，无指令发射开销）；数据流架构 RDU（SambaNova SN40L：1,040 个 PCU/PMU、638 TFLOPS bf16、520 MB 片上 SRAM、64 GB HBM、1.5 TB DDR，用 mesh 开关 + AGCU 做空间调度）；以及片上 SRAM / LPDDR 之间相差一到两个数量级的访问能耗（26 pJ vs 1200 pJ per 64 bit）。
  • 软件侧：把问题映射到机器的四步法——分解（decomposition）→ 分配（assignment）→ 编排（orchestration）→ 映射到硬件（mapping）；SPMD / SIMD 两层模型（ISPC 的 gang + task）；foreach 的四种实现（单实例、交错 interleaved、分块 blocked、动态原子分配）；metapipelining（流水线的流水线）+ 双缓冲 + token 控制的无锁数据流；以及自动调优 / 代价模型（分块尺寸、线程数、流水级数、并行维度的选择）。

• 在并行计算知识体系中的角色：这是全课程里唯一一讲把课程前面所有的”性能机制”倒过来当被优化的对象来看的讲次。前面我们学 Roofline、算术强度、SIMD、缓存分块、屏障与锁、reduce、work-span，是为了人去写出并行程序；本讲把这一整套分析变成学习型策略的代价模型与评价指标：无论是”学习型索引要不要多查一层模型”“预取器该不该在每次 L2 miss 上跑一个小 MLP”“32 卡时 allreduce 要不要和 GEMM 重叠”，答案都要用 work-span、算术强度、带宽/延迟公式和能耗预算算出来。同时它也是课程后半段（parallel deep learning、AI datacenter、heterogeneity / specialization）的收口：AI 系统的硬件给”用 AI 设计系统”提供了算力底座与片上存储，而 AI 系统的软件栈本身又充满了等待被学习型策略替换的启发式（分块尺寸、融合决策、并行维度、DRAM 调度）。

• 配套材料：
  • 本讲（CMU 15-418/15-618, Fall 2026, Oct 28, “AI in System Design”）在 Fall 2026 尚未发布讲义：主题是 Fall 2026 新增的，https://www.cs.cmu.edu/~418/lectures/ 之下没有对应的公开 PDF；历史学期在 /afs/cs/academic/class/15418-*/public/ 之下的讲义需要 CMU 登录，属未公开。因此本笔记不引用任何 Fall 2026 的幻灯片编号，讲义侧的事实基础全部来自下面两份已公开的补充材料与公开文献。
  • cs149_supp/aidatacenter.txt（已公开，Stanford CS149 Fall 2025 Lecture 12: Mapping AI Applications to the AI Datacenter，共 72 页）。它是本讲的硬件底座：HBM / 3D 堆叠 / TSV（slide 5、68）、HBM 带宽演进（Fury 512 GB/s → P100 720 GB/s → H100 3.2 TB/s，slide 70）、DRAM 阵列 / row buffer / precharge / RAS / CAS / burst / bank（slide 50–57）、内存控制器就是请求调度器且默认策略是 FR-FCFS（slide 63）、DDR4-2400 双通道 38.4 GB/s 与 ~13 ns CAS（slide 65）、能耗数字（1 pJ / 20 pJ / 26 pJ / 1200 pJ，slide 46；0.9 pJ / 5 pJ / 640 pJ，slide 47）、脉动阵列与数据流 RDU（slide 11–13）、metapipelining 与 kernel fusion（slide 16–26）、AI 模型里的并行维度 DP/TP/PP/EP/SP/CP 及通信原语 AllReduce / ReduceScatter / AllGather / All-to-All（slide 33–38）、计算-通信重叠的量化结果（8/16/32 socket 的 88.5%/77%/52% vs 实测 72%/75%/79%，slide 41）。
  • cs149_supp/thoughtprocess.txt（已公开，Stanford CS149 Fall 2025 Lecture 4: Parallelizing Code: The Programming Thought Process，共 74 页）。它是本讲的软件决策空间：创建并行程序的四步法（slide 30）、Amdahl 定律与 Summit 的 148,635,648 个 ALU（slide 32、36–37）、ISPC 的 SPMD 抽象与 SIMD 实现（slide 3、7、20–23）、foreach 的四种实现（slide 13）、interleaved 用 vmovaps 打包载入 vs blocked 用 vgatherdps gather（slide 10–11）、跨实例原语 reduce_add / broadcast / rotate（slide 18–19）、静态分配（C++11 std::thread）与动态任务分配（ISPC launch[100] + worker 线程池，slide 40–42）、red-black 染色重排与屏障/锁的成本（slide 50–72）。
  • 关于讲义首页写着 “Stanford CS149, Fall 2025”：这是补充材料的来源标注，不是错误；两份补充材料都是 CS149 的公开讲义，与本课 Fall 2026 的讲次编号无对应关系。
  • 未发布 / 需登录：本讲（Lecture 24）讲课录像在 Fall 2026 日程表中被注释隐藏，属未发布；Ed 讨论区、Autolab、Canvas 均需登录。同为 Fall 2026 尚未发布讲义的还有 Performance Analysis/Profiling、Transactional Memory 等。
  • 课程语境：Fall 2026 授课教师为 Brian Railing 与 Dimitrios Skarlatos；课程由 Kayvon Fatahalian 创建。课程主页 https://www.cs.cmu.edu/~418/，日程表 https://www.cs.cmu.edu/~418/schedule.html。
  • 本笔记引用的公开文献（属成熟公开知识，非课程特定内容）：Kraska et al., The Case for Learned Index Structures, SIGMOD 2018；Marcus et al., Benchmarking Learned Indexes, VLDB 2021；Ferragina & Vinciguerra, The PGM-index, PVLDB 2020；Marcus et al., Bao: Making Learned Query Optimization Practical, SIGMOD 2021；Hashemi et al., Learning Memory Access Patterns, ICML 2018 与 Voyager, MICRO 2021；Han et al., EIE, ICLR 2016（能耗数字来源，讲义 slide 47 同时引用）。
• 材料的公开状态一览（务必分清”本讲”与”补充材料”）

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

2.1 什么是 “AI in System Design”：两个方向与一个闭环

• 定义与目的：这门课前面讲的都是 Systems for AI——用专用硬件与编译器把 AI 跑快（脉动阵列、HBM、数据流 RDU、kernel fusion、metapipelining）。本讲的 “AI in System Design” 是反方向：用 AI/ML 去设计系统，即把系统内部那些”靠经验和拍脑袋定下来的规则”换成从数据里学出来的函数。系统设计的四步法（thoughtprocess.txt slide 30：Problem → Subproblems → Parallel Threads → Program → 机器）每一步都含有一批”配置选择”，而这些选择今天大多由手写启发式或静态常量给出：

  决策层级	手写启发式的例子	从哪份材料里看到	学习型替代
  数据表示 / 访问路径	B+-tree 的多级页查找；线性扫描	公开文献（learned index）	learned index（RMI / PGM）
  DRAM 请求排序	FR-FCFS（first-ready first-come-first-serve）	aidatacenter slide 63	学习型内存调度器
  缓存 / 预取	固定深度 stream / stride 预取器；LRU	aidatacenter slide 63 附近	学习型预取器 / 学习型替换
  分块与融合	固定的 tile size、固定的 kernel fusion 边界	aidatacenter slide 14、18–20	学习型代价模型 + 自动调优
  分配（assignment）	foreach 默认 interleaved；静态 blocked	thoughtprocess slide 13	由负载特征选择分配方案
  并行维度 / 流水	DP/TP/PP/EP 的度数靠人工网格搜索	aidatacenter slide 37、44	学习型并行策略搜索
  计算-通信重叠	“把 allreduce 放到后台线程”	aidatacenter slide 26、39–41	学习型流水调度

• 直观解释（”它是什么？”）：
  • 类比 1（系统里的”老师傅”）：一个跑了二十年的老机房管理员，凭直觉决定”哪个请求先发给内存”“这个循环切多大块”。他经验丰富但不会因为换了新机器而更新直觉（H100 的 roofline ridge point 是 CPU 的几十倍，老师傅的常数就全错了）。ML for systems 做的就是把老师傅的直觉换成一本会自己更新的手册：手册的输入是当前工作负载的特征（访存轨迹、队列长度、矩阵形状），输出是最优配置。
  • 类比 2（离线地图 vs 实时导航）：自动调优 / learned cost model 像离线地图——编译期或部署前把每个路口（每种配置）都跑一遍，代价高但决策慢一点没关系；学习型预取 / 学习型 DRAM 调度像实时导航——每一个路口（每一次访存）都要做一次决策，每次决策的预算只有纳秒和零点几纳焦。本讲最重要的工程洞察就藏在这个对比里：同一个模型，放在不同的决策点上，可行与不可行的差别是 6–9 个数量级。
• 图 1：ML for Systems 的在线闭环（软件执行模型）

 ┌────────────────────── AI in System Design：把启发式放进反馈回路 ──────────────────────┐
 │                                                                                       │
 │   ① 观测 (features)          ② 模型 (policy / cost model)      ③ 动作 (decision)      │
 │  ┌────────────────────┐      ┌──────────────────────────┐     ┌────────────────────┐ │
 │  │ 性能计数器 / 访存   │      │ 查表 / 线性模型 / 决策树  │     │ 预取地址、替换受害行│ │
 │  │ 轨迹 / 队列长度 /   │─────▶│ / GBDT / 小 MLP          │────▶│ tile size、线程数、 │ │
 │  │ PC、地址低位、shape │      │ θ 由离线训练或在线更新    │     │ 流水级数、路由、顺序│ │
 │  └────────────────────┘      └──────────────────────────┘     └─────────┬──────────┘ │
 │            ▲                                                            │            │
 │            │            ④ 反馈：重新训练 / 在线微调 / bandit 探索        ▼            │
 │  ┌─────────┴───────────┐                                    ┌────────────────────┐  │
 │  │ 目标系统 (被优化对象)│◀───────────────────────────────────│ 执行：cache / DRAM  │  │
 │  │ CPU/GPU/RDU、编译器、│        实测性能 / 能耗 / 尾延迟     │ 调度器、编译器、   │  │
 │  │ 索引、数据库、调度器 │───────────────────────────────────▶│ 运行时、数据库     │  │
 │  └─────────────────────┘                                    └────────────────────┘  │
 │                                                                                       │
 │  ⚠ 预算约束（本讲的核心张力）：                                                       │
 │     决策频率 × 单次推理成本 ≤ 该决策点能省的能耗/时间                                 │
 │     例：每 10 ns 一次的 DRAM 调度 ⇒ 单次推理预算 < 0.1 nJ  ⇒ 只能用几十个参数的查表     │
 │         每次编译一次 ⇒ 预算分钟级 ⇒ 可以跑 10^5 次真实测量 + 高斯过程                │
 └───────────────────────────────────────────────────────────────────────────────────────┘

• 性能特征与关键操作：这个闭环里，模型推理本身要在被它优化的同一台机器上跑。因此它带来三项新的系统开销，必须逐项量化：(a) 推理延迟（是否落在关键路径上？落在关键路径上就会直接变成串行部分，受 Amdahl 限制）；(b) 推理能耗（用 pJ/op 与 pJ/byte 换算，见 §4.5）；(c) 模型状态的空间与带宽（模型表多大？每决策触碰多少 cache line？）。一条经验法则：决策频率每提高 10 倍，可用模型参数大约减少 100 倍（因为预算按频率线性收紧，而可用的模型精度随参数增加呈对数增长）。

2.2 Learned Index Structures（学习型索引）

• 定义与目的：排序数组上的查找本质是 CDF（累积分布函数）求逆——pos = F(key)，其中 F 是”有多少个 key 小于等于 key”的单调函数。B+-tree 用 log₂ 次指针追逐去逼近这个函数，而学习型索引直接用模型拟合 F：顶层一个小模型给出粗略位置，下一层模型修正误差，最后只在误差界（error bound）ε 划定的窗口内做二分。它解决的问题是：索引本身的体积与指针追逐次数——索引从 GB 级降到 MB 级，每次查找的依赖访存从”树高”降到”模型层数 + 窗口内命中”。

• 直观解释（”它是什么？”）：
  • 类比（电话簿的页码估算）：要从一本按姓氏排序的电话簿里找 “Wang”，B+-tree 的做法的”二分查找”——每次都从正中间翻开，反复砍半；学习型索引的做法是先看姓的首字母在整本书里的位置分布（模型）：”W 大概在全书的 87% 处”，直接翻到 87% 附近，然后在前后几页里翻。如果这个分布是平稳的（模型准），你几乎一次就翻到；如果姓名分布变了（大量新增用户集中于 A 姓），模型的估计就会偏，你就得扩大搜索窗口——这就是学习型索引的脆弱点：它把”最坏情况有保证的 log n”换成了”平均情况很快，但依赖数据分布稳定”。
  • 另一个必要的类比（索引的”占地面积”）：B+-tree 像在图书馆里建了一整套目录卡片柜（每本书都要留位置），学习型索引像只留一张”每类书大致在第几排”的导览图——导览图当然不能告诉你第几本，但只要它把范围缩到几排之内，走过去翻一翻就够了。导览图不占地方，这就是学习型索引最稳的收益来源：空间缩小一到三个数量级，从而改善缓存行为。
• 图 2：两级 RMI（Recursive Model Index）的结构与查找路径（硬件/数据结构的存储布局）

                    两级 RMI：模型即索引，误差界 ε 决定搜索窗
  排序数组 K[0..n-1]（n = 2^28 个 64-bit key，位置 = 记录号，数组本身 2 GB）
  ┌─────────────────────────────────────────────────────────────────────────────────┐
  │ K[0]        K[n/4]        K[n/2]        K[3n/4]        K[n-1]                    │
  └─────────────────────────────────────────────────────────────────────────────────┘
        │
        │  ① 选一级段：在 b1[]（512 个"段首 key"，4 KB，L1 常驻）里二分 → 9 次 L1 访问
        ▼
  ┌─────────────────────────────────────────────────────────────────────────────────┐
  │ 一级（每段覆盖 n/512 = 2^19 个位置）                                             │
  │  b1[]  ┌────────┬────────┬────────┬─────┬────────┐        m1[] 每段 1 个模型    │
  │        │ key 段0│ key 段1│ key 段2│ ... │ key 段511│      pos ≈ a·key + b        │
  │        └────────┴────────┴────────┴─────┴────────┘                              │
  └─────────────────────────────────────────────────────────────────────────────────┘
        │
        │  ② 在该段的 256 个子段里，用 b2[]（2 KB，L1 常驻）二分 → 8 次 L1 访问
        ▼
  ┌─────────────────────────────────────────────────────────────────────────────────┐
  │ 二级（每个子段覆盖 2^19/256 = 2048 个位置）                                      │
  │  b2[512×256] (1 MB)   m2[512×256] (2 MB, 每模型 16 B)   e2[512×256] (0.5 MB)    │
  │  → 预测 pos2，并取出该子段的实测误差界 e2[c]                                     │
  └─────────────────────────────────────────────────────────────────────────────────┘
        │
        │  ③ 只在 [pos2 - e2, pos2 + e2] 内二分（由构造保证答案在其中）
        ▼
  ┌───────────────────────────────┐
  │ K[pos2-e2] … key … K[pos2+e2] │   窗口通常只有几个元素 → 1 次 cache line 命中
  └───────────────────────────────┘
   总索引体积 ≈ 4 KB + 1 MB + 2 MB + 0.5 MB ≈ 3.5 MB（对比 B+-tree 的 GB 级）
   每次查找的依赖 DRAM 访问 ≈ 1–2 次（对比朴素二分的 log2(2^28) = 28 次）

• 关键操作与性能特征：
  • 构造（build）：对每个段做一次线性插值拟合（O(1)），再实测该段的误差界（O(段长)）。总 work = Θ(n)（每位被访问常数次），而 B+-tree 的 bulk load 需要先排序再逐页填，work = Θ(n log n)（或 Θ(n) 但常数大得多）。Span = Θ(n/M1)（每个一级段可并行构造），并行度 = M1 = 512 起。
  • 查找（lookup）：延迟 = 一级边界二分（L1，~4 个周期/步）+ 二级边界二分（L1）+ 模型表访问（3.5 MB → L2/L3，约 15–40 ns）+ 窗口内二分（1 次 DRAM 缺失，约 90 ns）。总延迟约 130–200 ns；朴素 std::lower_bound 在 2 GB 数组上是 28 次依赖缺失 ≈ 2.5 µs。
  • 带宽：每次查找的 DRAM 流量 ≈ 2–3 个 64 B cache line ≈ 160 B；朴素二分每次 ≈ 28 × 64 B = 1.8 KB（相差 11 倍）。批量 10⁷ 次查找 ⇒ RMI ≈ 1.6 GB vs 朴素二分 ≈ 18 GB——这就是”学习型索引的价值最终体现为带宽”的含义。
• 表 1：B+-tree 与学习型索引（RMI / PGM）的定量对比（假设 n = 2²⁸ 个 64-bit key，cache line 64 B，DRAM 缺失延迟 90 ns，双通道带宽 38.4 GB/s）

必须知道的反面证据：Benchmarking Learned Indexes（VLDB 2021）这类系统性复现研究发现，在公平对比（同样的缓存占用、同样的并发控制、同样的写负载）下，学习型索引相对 B+-tree 的优势常常远小于最初论文报告的 1.5–3×，而且很大一部分收益来自”索引体积变小导致缓存行为改善”，而不是”模型预测更准”。这正是 “AI in System Design” 这一讲最该学到的态度：换掉一个启发式之前，先把整个系统的账算清楚——否则你只是在给瓶颈搬了个家。

2.3 Learned Cost Models 与 Autotuning（学习型代价模型与自动调优）

• 定义与目的：并行程序里有一大类决策，其特征是”离线可测、在线可查“：分块尺寸（tile size）、线程数、向量宽度、循环融合的边界、数据布局（AoS vs SoA）、流水级数、DP/TP/PP 的度数、通信-计算重叠的方案。这些决策的搜索空间巨大（aidatacenter slide 14 把数据流编译器的流水线写成”Tiling → Parallelization → Metapipelining → Place & Route → Codegen”，每一级都是一个搜索维度），而手写常量在每次硬件换代时就失效。learned cost model 的目的就是：用少量真实测量拟合一个代理模型（surrogate），再用它去预测没测过的配置的性能，从而把”穷举”压成”少量测量 + 大量预测”。

• 直观解释（”它是什么？”）：
  • 类比（请人试鞋）：穷举搜索是把所有尺码的鞋都试一遍（慢但准）；解析代价模型是拿公式算（快但常常算不准，比如 roofline 完全忽略了 bank 冲突和调度开销）；学习型代价模型是先试 6–8 双，然后根据”鞋码 ↔ 合脚程度”的关系推断出你的尺码，再直接推荐一双——最后还要再试一下验证（train/test）。本节 §3 的第二个代码示例就是这个过程的完整实现：在 6 个 tile size 上测量 → 拟合模型 → 预测未见过的矩阵尺寸上的最优 tile → 验证。
  • 另一个类比（预报 vs 实测）：解析 roofline 是”气象模型”（物理正确但忽略很多细节），学习型代价模型是”统计回归”（不解释机理，但在训练分布内很准）。两者必须一起用：roofline 给出上界（你去不了的地方），学习型模型给出可行域内的最优选择（你实际能到的地方）。
• 图 3：自动调优 / 学习型代价模型的闭环与搜索空间（软件执行模型）

   ┌──────────────┐  配置 x=(T, threads, order, fusion)   ┌──────────────────────┐
   │ 搜索器        │──────────────────────────────────────▶│  目标系统/内核       │
   │ · 网格 / 随机 │                                       │  GEMM 分块、卷积融合 │
   │ · 贝叶斯优化  │◀──────────────────────────────────────│  流水级数、并行维度  │
   │ · bandit/进化 │   测量 y = (ms, GFLOPS, J, tail)      └──────────────────────┘
   └──────┬───────┘                                                 │
          │ 预测 ŷ(x; θ)                                             │ 真实执行
          ▼                                                          ▼
   ┌────────────────────────────────────────┐              ┌──────────────────────┐
   │ 代价模型 cost model                    │   训练 θ     │ 采样点 (x_i, y_i)    │
   │  ① 解析式（roofline 上界）:            │◀─────────────│  只需 6–20 个点      │
   │     t = max(2MNK/Peak, traffic/BW)     │              └──────────────────────┘
   │  ② 效率修正项:                         │
   │     eff(x) = 1/(1 + c₁/T + c₂·T²/n²)   │   ← 拟合两个物理意义明确的常数：
   │  ③ 学习式（可选）: GBDT / 小 MLP        │      ① 每块固定开销 ∝ 1/T
   │     特征 = shape、cache 大小、SM 数     │      ② 块数不足 ⇒ 并行度不足 ∝ T²/n²
   └────────────────────────────────────────┘
   迁移学习：把在机器 A 上拟合的 θ 作为机器 B 的初值，只需 1–2 次测量就能校准

• 关键操作与性能特征：
  • 搜索空间的维度爆炸：仅”分块 + 循环顺序 + 向量化 + 线程数 + 融合边界”就轻松超过 10⁶ 个配置点，而每个配置的一次真实测量往往是几十毫秒到几秒。穷举的 work = |配置空间| × 单次测量时间，这在工程上不可接受，于是 Work 必须从”全空间”降到”O(log)|空间|“（贝叶斯优化）或”O(1) 预测”（拟合后）。
  • 训练数据的成本：每次测量本身就是一次完整的程序运行（含编译、加载、预热）。这与 §2.2 的索引构造不同：索引构造的 work 是 Θ(n) 的一次性代价，而调优的 work 是 |采样| × 内核运行时间，并且必须在新 shape/新硬件上重新支付。
  • 泛化与分布漂移：模型在 shape 分布内的插值通常很准（这也是为什么用 log T 而不是 T 做特征），但外推（更大的矩阵、更多的线程、更强的 GPU）常常失效，必须保留”不确定时回到安全默认值”的兜底路径。

2.4 学习型内存系统策略：预取、替换与 DRAM 调度

• 定义与目的：内存系统的每一个决策都是一个”顺序 + 选择”问题：内存控制器先服务哪个请求（aidatacenter slide 63：默认策略是 FR-FCFS——先服务”当前行已打开”的请求以最大化行局部性，其余按 FIFO）、预取哪个地址、淘汰哪一行（LRU）。这些启发式在单一负载下不错，但在混合负载（一个流式扫描 + 一个随机访问）下会互相伤害。学习型策略的输入是控制器天然拥有的丰富状态（每个 bank 的队列、每行的 open/close 状态、请求的 PC、地址低位、队列长度），输出是一个更优的排序或地址预测。

• 直观解释（”它是什么？”）：
  • 类比（超市收银台的排队规则）：FR-FCFS 就是”先把购物车里东西已经摆在收银台上的顾客结掉，其余的按到达顺序”——它在大多数时候很合理，但如果那个”东西摆好”的顾客要结 500 件商品，后面只买一瓶水的顾客就被饿死了。学习型调度则像一位会看队伍构成的领班：根据”每个人的商品数、是否已摆好、通道拥堵情况”动态决定先开哪个收银台——更优，但领班自己也要发工资（推理开销），而且如果他的判断依据是过时的培训内容（分布漂移），可能比死板的规则更糟。
  • 关键约束（DRAM 侧）：aidatacenter slide 54 强调 data pins 是最稀缺的资源，只在很小一部分时间里被占用；slide 56 说明 多个 bank 允许请求流水化（一个 bank 做 precharge/activate，另一个 bank 在传输数据），slide 55 说明 burst mode 把延迟摊到更大的传输上。所以 DRAM 调度策略的目标函数就是”让数据总线始终忙 + 让行命中率尽量高”——这是一个非常明确、非常适合用一个小模型去学的问题，也是学习型策略在硬件里最有希望的落点之一。
• 图 4：内存控制器（请求调度器）+ 学习型仲裁（硬件结构）

   内存控制器 = 请求调度器（讲义 slide 63）：输入是 LLC 的缺失请求，输出是 DRAM 命令序列
                               ┌────────────────────────────────────────┐
   LLC (L3) 缺失 ──┬──▶ ┌──────────────┐   ┌───────────────┐   ┌──────────────┐
                   ├──▶ │ bank0 请求队列│──▶│ 行缓冲 row B0 │   │  策略 policy │
                   ├──▶ │ bank1 请求队列│──▶│ 行缓冲 row B1 │◀──│ ① FR-FCFS    │
                   ├──▶ │ bank2 请求队列│──▶│ 行缓冲 row B2 │   │  (基线启发式)│
                   └──▶ │ bank3 请求队列│──▶│ 行缓冲 row B3 │   │ ② 学习型排序 │
                        └──────────────┘   └───────┬───────┘   │   输入特征：  │
                               ▲                   │           │   队列长度、  │
                               │                   │           │   行命中、PC、│
                        请求可合并？               ▼           │   地址低位、  │
                        (burst mode 的基础)   数据引脚 (8 bit)  │   读写比例    │
                                              ═══════════════  │ ③ 小型线性/  │
                                              64-bit 内存总线  │    查表模型   │
                                                              └──────────────┘
   时间轴：RAS(row activate) → CAS(column access) → burst 传输 → PRE(precharge)
   ┌────┬─────────┬─────────┬─────────┬────┬─────────┬─────────┐
   │PRE │  RAS    │  CAS    │ burst×N │PRE │  RAS    │  CAS    │   ← 总线空闲的坑
   └────┴─────────┴─────────┴─────────┴────┴─────────┴─────────┘
        └── ~10 ns ─┴── ~10 ns ─┴─ 传输 ─┘
   调度器要做的就是"用别的 bank 的命令把这个坑填上"（bank 级流水），
   而"填哪个坑最划算"正是学习型策略要预测的量。

• 关键操作与性能特征（带宽、延迟、吞吐）：
  • 延迟：从”行已打开”到数据上总线的 CAS ≈ 10–13 ns（slide 51、65）；从”行未打开”则要 PRE + RAS + CAS ≈ 30 ns。同一个 DRAM 的延迟差别达 3 倍——这解释了为什么调度器（而不是更快的 DRAM）是主要优化对象。
  • 带宽 vs 延迟必须一起算（Little’s law）：要维持带宽 BW，在途字节数必须 ≥ BW × 延迟。DDR4-2400 单通道 19.2 GB/s × 13 ns ≈ 250 B ≈ 4 个 cache line；H100 的 3.2 TB/s × 500 ns ≈ 1.6 MB（全芯片在途）≈ 每 SM 约 12 KB ≈ 190 个 cache line。学习型预取/调度的第一性约束不是”预测多准”，而是”能不能造出足够的在途请求”——再准的预取器，若队列深度不够也填不满带宽。
  • 吞吐：aidatacenter slide 65 的 DDR4-2400 数字可以直接算：64-bit × 1.2 GHz × 2（双沿）× 2 通道 = 38.4 GB/s。

2.5 硬件底座：学习型策略跑在什么机器上

• 定义与目的：学习型策略的可行性完全由它运行的那台机器决定：推理要多快、模型能占多少 SRAM、一次访存多少 pJ。aidatacenter.txt 给出的正是这块底座。它解决的是”在哪儿算、数据从哪儿来“的问题：HBM + 3D 堆叠把带宽从 DRAM 的几百 GB/s 拉到几 TB/s；脉动阵列把稠密矩阵乘的单位能耗压到最低；数据流 RDU 用片上大容量 SRAM（520 MB vs H100 的约 100 MB）+ metapipelining 把中间结果留在片上。

• 直观解释（”它是什么？”）：
  • 类比（把水库修到工厂隔壁）：传统 DRAM 像城市边缘的水库——水很便宜，但管道又长又细（64-bit 内存总线），每次取水都要等水压上来（延迟）。HBM 通过 3D 堆叠 + silicon interposer 把水库修到了车间地板下面（1024-bit/stack 的宽接口，走线只有毫米级），于是管道变宽、距离变短、单位能耗下降（slide 6：更高带宽、更高能效、更小封装）。代价是容量与 3D 堆叠的制造难度。
  • 类比（流水线的工人 vs 会传话的工人）：脉动阵列像一条传送带上的工人：每个人只做一次乘加，然后把结果递给下一个人（数据在相邻 PE 间只走一个寄存器到寄存器的距离），没人需要跑仓库取料（无指令发射、无寄存器文件读写的开销）。数据流架构 RDU 像一群按图纸自动流转的车间：数据带着”token”在 PCU/PMU 之间流动，下游知道该什么时候干活（token 控制），因此不需要锁（slide 21：Dataflow execution with token control ⇒ no lock-based synchronization）。
• 图 5：HBM 与 3D 堆叠（TSV + interposer，硬件结构；依据 slide 5/68）

                     3D 堆叠 + TSV + silicon interposer
        ┌─────────────────┐
        │   DRAM die 3    │
        ├─────────────────┤  TSV (through-silicon via)：贯穿各 die 的垂直互连，
        │   DRAM die 2    │  把逻辑层与 DRAM 之间做成"高度并行的连接"
        ├─────────────────┤  （不是传统的一根总线，而是成千上万个垂直通孔）
        │   DRAM die 1    │
        ├─────────────────┤
        │  logic die      │  ← 集成 memory controller：管理来自处理器的请求
        │  (控制器/基底层)│
        └────────┬────────┘
                 │  1024-bit 接口 / stack（H100: 6 stack × 1024 bit = 6144 bit）
     ════════════╧═══════════════════════════════════════════
       silicon interposer：高带宽、短距离互连（"把内存搬到处理器旁边"）
     ════════════╤═══════════════════════════════════════════
                 │
        ┌────────┴────────┐
        │   Processor     │   CPU: 64-bit 内存总线 → DRAM
        │  (GPU / SM 阵列)│   GPU: 4096/6144-bit → HBM stack
        └─────────────────┘
     带宽演进（slide 70）：AMD Fury 2015  4×1024-bit → 512 GB/s
                          NVIDIA P100 2016 4×HBM2 → 720 GB/s
                          NVIDIA H100 2022 6×HBM3 → 3.2 TB/s（80 GB）

• 图 6：脉动阵列（systolic array，硬件结构）

   稠密矩阵乘 C[M,N] += A[M,K]·B[K,N]：权重 B 预载并"静止"在阵列里，A 从左边界逐拍注入
        ┌─────┬─────┬─────┬─────┐
  a0 ──▶│ PE  │────▶│ PE  │────▶│ PE  │────▶│ PE  │──▶   (a 向右脉动)
        ├─────┼─────┼─────┼─────┤
  a1 ──▶│ PE  │────▶│ PE  │────▶│ PE  │────▶│ PE  │──▶
        ├─────┼─────┼─────┼─────┤
  a2 ──▶│ PE  │────▶│ PE  │────▶│ PE  │────▶│ PE  │──▶
        ├─────┼─────┼─────┼─────┤
  a3 ──▶│ PE  │────▶│ PE  │────▶│ PE  │────▶│ PE  │──▶
        └──┬──┴──┬──┴──┬──┴──┬──┘
           │     │     │     │
           ▼     ▼     ▼     ▼
          c0    c1    c2    c3      (部分和向下脉动，在对角线时刻依次出结果)
   每个 PE：1 个乘法器 + 1 个累加器 + 1 个寄存器（没有指令、没有地址计算、没有 cache）
   ⇒ 面积极小、能耗极低（数据只在相邻 PE 间移动，几乎不经过长导线）
   ⇒ 代价：只有"能整齐映射到阵列尺寸上的稠密矩阵乘"才能高效运行；
     稀疏、不规则、动态 shape 的问题映射效率骤降 ⇒ 这正是"学习型代价模型/调度"
     要处理的一类问题（该不该把算子交给脉动阵列、要不要 pad、要不要换布局）。

• 图 7：RDU 上的 metapipelining 与 token 控制的无锁数据流（硬件 + 软件执行模型；依据 slide 13/16/21/24）

  SN40L RDU：1,040 个 PCU/PMU，638 TFLOPS bf16，520 MB 片上 SRAM，64 GB HBM，1.5 TB DDR
  ┌──────────────────────────── 片上 mesh（S = switch，高带宽互连）──────────────────────┐
  │  ┌──────────┐   ┌──────────┐   ┌──────────┐   ┌──────────┐   ┌──────────┐            │
  │  │  AGCU    │──▶│  PMU     │──▶│  PCU     │──▶│  PMU     │──▶│  AGCU    │──▶ 片外    │
  │  │ 地址生成 │   │ a_tile   │   │ MAT_MUL  │   │ c_tile   │   │          │   HBM/DDR │
  │  │ 合并单元 │   │ 0.5 MB   │   │ 16×8 bf16│   │ 双缓冲   │   │          │            │
  │  │ (片外门户)│   │ 高灵活寻址│   │ systolic │   │          │   │          │            │
  │  └──────────┘   └──────────┘   └──────────┘   └──────────┘   └──────────┘            │
  │        ▲              ▲              ▲              ▲                                 │
  │        │              │              │              │  级间以 token/credit 握手        │
  │  ══════╧══════════════╧══════════════╧══════════════╧═══  ⇒ 无锁同步（slide 21）      │
  │  METAPIPE(M/MM){ a=LOAD_TILE(A) ; METAPIPE(N/NN){ b=LOAD_TILE(B) ;                    │
  │                          c=MAT_MUL(a,b,row_par=4) ; c_t=BUFFER(c) ; STORE_TILE(C,c_t)}}│
  └───────────────────────────────────────────────────────────────────────────────────────┘
   时间轴（4 级 metapipeline；双缓冲让 load 与 compute 重叠，稳态后每拍出一个 tile）：
     t :   0    1    2    3    4    5    6    7    8
   ld  : │ T0 │ T1 │ T2 │ T3 │ T4 │ T5 │ T6 │ T7 │ T8 │   AGCU 持续把 tile 从 HBM 拉进来
   mm  : │    │ T0 │ T1 │ T2 │ T3 │ T4 │ T5 │ T6 │ T7 │   PCU 只在数据就绪时才动（token）
   buf : │    │    │ T0 │ T1 │ T2 │ T3 │ T4 │ T5 │ T6 │   PMU 双缓冲吸收级间速率差
   st  : │    │    │    │ T0 │ T1 │ T2 │ T3 │ T4 │ T5 │   写回不占用 HBM 带宽的关键：融合
   ⇒ 级间速率不匹配（imbalanced stages）由双缓冲吸收；中间结果不出片
   ⇒ 整个 decoder 融合成 1 个 kernel：每 token 的 kernel 调用从 GPU 的 ~800 次降到 3 次

• 表 2：访问与运算的开销（数据来自 aidatacenter.txt slide 46/47，两套来源同时列出；本讲的能量预算全靠这张表）

两个口径相差 20 倍（FP op 0.9 pJ vs 20 pJ）——这不矛盾：slide 46 是较先进工艺 + 只计逻辑运算本身的功耗，slide 47 是 45 nm 估算且常被用来做”数据搬运 vs 计算”的相对比较。做能耗预算时必须写明用的是哪一套数字，否则结论会差一个数量级以上（§4.5 会看到这个差别直接决定”学习型预取器是否可行”）。

2.6 软件执行模型：从 foreach 的分配策略到 metapipelining

• 定义与目的：thoughtprocess.txt 把”写一个并行程序”拆成四步：分解（找独立工作）→ 分配（把工作给 worker）→ 编排（通信/同步/数据布局/调度）→ 映射到硬件。这四步里，“分配”和”编排”正是学习型策略可以接管的部分：一个 ISPC 编译器/运行时可以在 interleaved / blocked / dynamic 之间选，一个数据流编译器可以在 tile 大小、流水级数、融合边界之间选。

• 直观解释（”它是什么？”）：
  • 类比（分糖果的三种分法）：interleaved（交错）是”一人一颗轮着发”——每个人手里的糖在记忆里是相邻的（对 SIMD 打包载入极友好，vmovaps 一条指令搞定 8 个）；blocked（分块）是”每人发一整段”——每个人手里的糖在记忆里隔得很远（SIMD 必须用 vgatherdps 一条条去抓，慢好几倍），但如果每颗糖的”加工时间”不均匀，分块反而更均衡。dynamic（动态）是”发一颗算一颗，算完再来拿”——最均衡，但每次拿糖都要排队（原子操作 + cache line 争用）。哪一种最优，取决于”加工时间的方差”和”访问的连续性”这两个特征——这正是学习型策略的输入。
  • 类比（流水线 vs 一次性做完）：metapipelining 是”把一条长流水线切成互相重叠的小段”（aidatacenter slide 16 称其为 “a pipeline of pipelines”）：它把循环体里的并行模式（map/reduce/zip）变成流水级，用双缓冲吸收各级速率差，让多个循环迭代同时”在路上”。融合（fusion）与 metapipelining 的分界本身就是个决策：能融合就融合（省 HBM 流量），不能融合（依赖太复杂）就用 metapipelining（aidatacenter slide 16 明确写了 “Metapipelining can work when fusion does not”）。
• 图 8：ISPC 的两层并行（gang × task）与 foreach 的分配策略（软件执行模型；依据 thoughtprocess slide 3/7/13/22/42）

   第一层：gang（SIMD，一个线程内）                第二层：task（多核，线程池）
   ┌──────────────────────────────────┐          ┌──────────────────────────────────┐
   │ export void f(...) {             │          │ void foo(...) {                  │
   │   foreach (i = 0 ... N) { ... }  │          │   launch[100] my_task(...);      │
   │ }                                │          │ }   ← 100 个 task 扔进任务表     │
   │ programCount = 8 (AVX2 fp32)     │          └────────────────┬─────────────────┘
   │ instance 0..7 各持 programIndex  │                           │
   └──────────────────────────────────┘                           ▼
   ┌───────────────────────────────────────────────────────────────────────────────┐
   │  worker 线程池：worker0  worker1  worker2  worker3                            │
   │       │        │        │        │                                            │
   │       └────────┴───┬────┴────────┘   任务队列: t0 t1 t2 t3 t4 … t99          │
   │                    │                  next-task 指针 ⇒ 谁空谁取（动态分配）    │
   └───────────────────────────────────────────────────────────────────────────────┘

   foreach 的四种实现（同一份代码，编译器/运行时可任选一种）：
   ┌─ impl 1：只让 instance 0 跑全部迭代      → 纯标量，无 SIMD 加速            ┐
   ├─ impl 2：interleaved  i = loop_i + programIndex                             │
   │     迭代: 0 1 2 3 4 5 6 7 | 8 9 …      访问: 8 个连续 float → vmovaps ✓    │
   ├─ impl 3：blocked      i = start + loop_i, start = programIndex*count         │
   │     迭代: 0 8 16 24 …  | 1 9 17 …      访问: 8 个跨步 float → vgatherdps ✗  │
   └─ impl 4：dynamic       i = atomic_add_local(&nextIter, 1)                    │
         迭代: 谁先取到谁算                 均衡最好，但每次取号要原子操作        ┘

• 关键操作与性能特征（延迟、带宽、吞吐）：
  • SIMD 载入方式决定访存吞吐：interleaved 的 vmovaps（打包载入）在 Skylake 级核心上约 0.5 周期/8 float；vgatherdps（gather）需要逐元素访问，吞吐差约一个数量级（具体数字随微架构而异）。所以 blocked 分配在”访问必须连续”这一条上要付 10 倍的载入代价，除非该负载本来就带宽受限（此时 gather 的代价会被内存墙掩盖——这又是一个”必须做 Roofline 才能判断”的例子）。
  • 动态分配的开销与 Amdahl：每次取号一次原子操作（约 20–100 ns 的争用成本）。若一个任务只有 1 µs 的有效工作，取号开销就占 2–10%；若任务有 100 µs 工作，开销可忽略。“任务粒度”因此是一个由 原子开销 / 任务有效工作 决定的、可以用学习型模型预测的决策。
  • 编排的代价（屏障与锁）：thoughtprocess slide 60/68 的网格求解器例子给出了最经典的反面教材：把 lock/unlock 放进内层循环（每次 (i,j) 更新都加锁）会让程序的并行度几乎归零；改成”先在线程内累加 myDiff，再在循环外锁一次”，锁开销从 O(N²) 次降到 O(1) 次（每线程每轮一次）。slide 72 又把它进一步优化成”用三个不同的 diff 变量轮流使用，把三个屏障降到一轮一个“——用空间（footprint）换依赖的消除。

3. 代码示例与性能分析

3.1 示例一：两级 RMI 学习型索引 vs 多级页式索引 vs 朴素二分（C++17 + OpenMP）

// ============================================================================
// learned_index.cpp — 学习型索引 (两级 RMI, 误差界驱动搜索窗) vs 多级页式索引 vs 朴素二分
// 编译: g++ -O3 -march=native -fopenmp -std=c++17 learned_index.cpp -o learned_index
// 运行: OMP_NUM_THREADS=8 ./learned_index 268435456 10000000
//        参数1 = key 个数 n（默认 1<<24 = 16.7M）
//        参数2 = 查询次数 nq（默认 4*10^6）
// 说明: -O3 + -march=native 打开向量化与本地指令集；这是 release 优化编译。
// ============================================================================
#include <algorithm>
#include <chrono>
#include <cmath>
#include <cstdint>
#include <cstdio>
#include <cstdlib>
#include <random>
#include <vector>
#include <omp.h>

static double now_ms() {
  using namespace std::chrono;
  return duration<double, std::milli>(steady_clock::now().time_since_epoch()).count();
}

// ---------- 数据：单调递增、密度分两段的 key（"热区 + 冷区"）----------
// 单条直线无法拟合整条 CDF，因此顶层模型必然有误差 —— 这正是需要第二层的原因。
static void make_keys(std::vector<uint64_t>& K) {
  const int64_t n = (int64_t)K.size();
  for (int64_t i = 0; i < n; ++i) {
    if (i < n / 2) K[i] = (uint64_t)i * 2;                          // 前半段：密度 1/2
    else           K[i] = (uint64_t)n + (uint64_t)(i - n / 2) * 8;  // 后半段：密度 1/8
  }
}

// ---------- 线性模型 pos ≈ a*key + b ----------
struct Model {
  double a = 0.0, b = 0.0;
  int64_t predict(uint64_t k) const { return (int64_t)(a * (double)k + b); }
};

// 用段的两端点做线性插值拟合：O(1)，适合构建上百万个模型
static Model fit_range(const std::vector<uint64_t>& K, int64_t lo, int64_t hi) {
  Model m;
  const double x0 = (double)K[lo], y0 = (double)lo;
  const double x1 = (double)K[hi], y1 = (double)hi;
  if (x1 <= x0) { m.a = 0.0; m.b = y0; return m; }
  m.a = (y1 - y0) / (x1 - x0);
  m.b = y0 - m.a * x0;
  return m;
}

// ---------- 方案 A：两级 RMI（学习型索引）----------
struct RMI {
  int M1 = 512, SUB = 256;                 // 一级 512 段，每段 256 个子段
  int64_t n = 0;
  std::vector<uint64_t> b1, b2;            // 边界 key：b1=512*8B=4KB, b2=512*256*8B=1MB
  std::vector<Model>    m1, m2;            // 模型：m1=8KB, m2=512*256*16B=2MB
  std::vector<int32_t>  e2;                // 每个二级子段的实测最大误差（0.5MB）
  int32_t max_e1 = 0;

  int64_t seg_lo(int i) const { return (int64_t)((double)n * i / M1); }
  int64_t seg_hi(int i) const { return (int64_t)((double)n * (i + 1) / M1) - 1; } // 闭区间

  void build(const std::vector<uint64_t>& K) {
    n = (int64_t)K.size();
    b1.resize(M1); m1.resize(M1);
    b2.resize((size_t)M1 * SUB); m2.resize((size_t)M1 * SUB); e2.resize((size_t)M1 * SUB);
    for (int i = 0; i < M1; ++i) {
      const int64_t lo = seg_lo(i), hi = seg_hi(i);
      b1[i] = K[lo];
      m1[i] = fit_range(K, lo, hi);
      // 一级误差（仅用于统计/诊断；本设计的正确性不依赖它）
      for (int64_t p = lo; p <= hi; ++p) {
        int64_t e = std::llabs(m1[i].predict(K[p]) - p);
        if (e > max_e1) max_e1 = (int32_t)e;
      }
      // 二级：把本段等分成 SUB 个子段，每个子段一个模型 + 实测误差界
      for (int t = 0; t < SUB; ++t) {
        const int64_t slo = lo + (hi - lo + 1) * t / SUB;
        const int64_t shi = lo + (hi - lo + 1) * (t + 1) / SUB - 1;
        const size_t idx = (size_t)i * SUB + t;
        b2[idx] = K[slo];
        m2[idx] = fit_range(K, slo, shi);
        int64_t e = 0;
        for (int64_t p = slo; p <= shi; ++p) {
          int64_t d = std::llabs(m2[idx].predict(K[p]) - p);
          if (d > e) e = d;
        }
        e2[idx] = (int32_t)e;
      }
    }
  }

  // 纯函数、可安全并发调用：这是查找路径的全部
  int64_t lookup(const std::vector<uint64_t>& K, uint64_t key, int64_t* fallbacks = nullptr) const {
    // ① 一级：在 512 个段首 key 里二分（4 KB，L1 常驻）
    int i = (int)(std::upper_bound(b1.begin(), b1.end(), key) - b1.begin()) - 1;
    if (i < 0) i = 0;
    // ② 二级：在本段的 256 个子段边界里二分（2 KB，L1 常驻）
    const uint64_t* base = b2.data() + (size_t)i * SUB;
    int c = (int)(std::upper_bound(base, base + SUB, key) - base) - 1;
    if (c < 0) c = 0;
    const size_t idx = (size_t)i * SUB + c;
    // ③ 用子段模型+误差界划定搜索窗（由构造保证真值在窗内）
    const int64_t seg_lo_p = (int64_t)((double)n * i / M1);
    const int64_t seg_hi_p = (int64_t)((double)n * (i + 1) / M1) - 1;
    int64_t q = m2[idx].predict(key);
    int64_t lo = std::max(seg_lo_p, q - e2[idx] - 1);
    int64_t hi = std::min(seg_hi_p + 1, q + e2[idx] + 2);
    if (lo < hi && K[lo] <= key && key <= K[hi - 1])
      return (int64_t)(std::lower_bound(K.begin() + lo, K.begin() + hi, key) - K.begin());
    if (fallbacks) ++(*fallbacks);           // 兜底：极端分布漂移时回退全局二分
    return (int64_t)(std::lower_bound(K.begin(), K.end(), key) - K.begin());
  }
};

// ---------- 方案 B：多级页式索引（B+-tree 风格的指针追逐）----------
struct PagedIndex {
  static constexpr int FANOUT = 256;      // 每页 256 个 key（8B key ⇒ 2 KB 页）
  const std::vector<uint64_t>* K = nullptr;
  std::vector<std::vector<uint64_t>> upper;   // 自底向上抽稀的"内部节点"数组
  void build(const std::vector<uint64_t>& keys) {
    K = &keys;
    upper.clear();
    std::vector<uint64_t> cur;
    for (size_t i = 0; i < keys.size(); i += FANOUT) cur.push_back(keys[i]);
    while (cur.size() > 1) {
      upper.push_back(cur);
      std::vector<uint64_t> nxt;
      for (size_t i = 0; i < cur.size(); i += FANOUT) nxt.push_back(cur[i]);
      cur.swap(nxt);
    }
    upper.push_back(cur);                 // 顶层（1 个元素），统一下降逻辑
  }
  int64_t lookup(uint64_t key) const {
    const std::vector<uint64_t>& L = *K;
    int64_t block = -1;                   // -1 = 整个数组
    for (int l = (int)upper.size() - 1; l >= 0; --l) {
      const std::vector<uint64_t>& a = upper[l];
      size_t lo = (block < 0) ? 0 : (size_t)block * FANOUT;
      size_t hi = (block < 0) ? a.size() : std::min(a.size(), lo + (size_t)FANOUT);
      size_t p = (size_t)(std::upper_bound(a.begin() + lo, a.begin() + hi, key) - a.begin());
      block = (int64_t)((p == lo) ? lo : p - 1);
    }
    size_t lo = (block < 0) ? 0 : (size_t)block * FANOUT;
    size_t hi = std::min(L.size(), lo + (size_t)FANOUT);
    return (int64_t)(std::lower_bound(L.begin() + lo, L.begin() + hi, key) - L.begin());
  }
};

// ---------- 驱动：正确性校验 + 并行批量计时 ----------
int main(int argc, char** argv) {
  const int64_t n  = (argc > 1) ? atoll(argv[1]) : (1LL << 24);
  const int64_t nq = (argc > 2) ? atoll(argv[2]) : 4000000;
  std::vector<uint64_t> K((size_t)n);
  make_keys(K);

  // 生成查询 key（可能不存在于数组中，用来验证 lower_bound 语义）
  std::vector<uint64_t> Q((size_t)nq);
  { std::mt19937_64 rng(12345); for (int64_t i = 0; i < nq; ++i) Q[i] = rng() % (2ULL * (uint64_t)n); }

  RMI rmi;      double t0 = now_ms(); rmi.build(K);      double tb_rmi = now_ms() - t0;
  PagedIndex pg; t0 = now_ms();       pg.build(K);       double tb_pg  = now_ms() - t0;
  std::printf("n=%lld  nq=%lld\n", (long long)n, (long long)nq);
  std::printf("构造: RMI %.1f ms | PagedIndex %.1f ms | RMI 索引体积 %.2f MB (B+-tree 风格 %.2f MB)\n",
              tb_rmi, tb_pg,
              (rmi.b1.size()*8.0 + rmi.b2.size()*8.0 + rmi.m1.size()*16.0 +
               rmi.m2.size()*16.0 + rmi.e2.size()*4.0) / 1048576.0,
              (double)n * 8.0 / 1048576.0);
  std::printf("RMI 一级最大误差 max_e1=%d（说明单条直线确实拟合不了整个 CDF）\n", rmi.max_e1);

  // 单线程校验：RMI / PagedIndex 必须与 std::lower_bound 完全一致
  int64_t bad = 0, fb = 0;
  for (int64_t i = 0; i < std::min<int64_t>(nq, 200000); ++i) {
    int64_t ref = (int64_t)(std::lower_bound(K.begin(), K.end(), Q[i]) - K.begin());
    if (rmi.lookup(K, Q[i], &fb) != ref) ++bad;
    if (pg.lookup(Q[i]) != ref) ++bad;
  }
  std::printf("校验: 不一致 %lld 次（必须为 0）；RMI 兜底回退 %lld / 200000 次\n",
              (long long)bad, (long long)fb);

  // 并行批量计时（静态分配；无共享写入 ⇒ 无数据竞争、无伪共享）
  auto bench = [&](const char* name, auto fn) {
    double t = now_ms();
    int64_t acc = 0;
    #pragma omp parallel for schedule(static) reduction(+:acc)
    for (int64_t i = 0; i < nq; ++i) acc += fn(Q[i]);
    const double ms = now_ms() - t;
    std::printf("%-22s %8.2f ms   %6.1f ns/查询   校验和 %lld\n",
                name, ms, ms * 1e6 / (double)nq, (long long)(acc & 1));
  };
  bench("朴素 lower_bound", [&](uint64_t k){ return (int64_t)(std::lower_bound(K.begin(), K.end(), k) - K.begin()); });
  bench("多级页式索引",     [&](uint64_t k){ return pg.lookup(k); });
  bench("学习型索引 RMI",   [&](uint64_t k){ return rmi.lookup(K, k); });
  return 0;
}

【代码做什么？】

1. make_keys 生成 单调递增但密度分两段的 key：前半段密度 1/2、后半段 1/8。这样单条直线无法拟合整条 CDF（程序会打印出 max_e1 远大于 0 来证明这一点），必须引入第二级模型——这正是 RMI 的存在理由。
2. RMI::build 分两级构造：一级把有序数组按位置等分成 512 段，每段用两端点做线性插值得到 pos ≈ a·key+b，并把段首 key 存进 b1（4 KB）；二级把每个一级段再等分成 256 个子段，各拟合一个模型，并且逐元素实测该子段的最大误差 e2[]——这是”误差界”的来源。构造总成本 Θ(n)（每位被访问常数次），且各段之间完全独立，可以并行。
3. RMI::lookup 是没有分支预测惊喜的确定性路径：在 512 个段首 key 里二分（4 KB，L1 命中）→ 在该段的 256 个子段边界里二分（2 KB，L1 命中）→ 用子段模型预测位置 → 在 [q-e2-1, q+e2+2] 这个由构造保证包含真值的窗口里 lower_bound。窗口通常只有几个元素（1 个 cache line）。若窗口校验失败（分布漂移、查询 key 落在段边界外），回退到全局二分并计数——学习型索引必须永远带一条兜底路径。
4. PagedIndex 是公平的对照组：它就是 B+-tree 的”每页 256 项、逐层抽稀”的指针追逐模型，查找的依赖访存级数 = 层数。
5. main 先做单线程正确性校验（RMI 与 PagedIndex 必须和 std::lower_bound 逐位一致），再对三种方案做并行批量计时。计时用 #pragma omp parallel for schedule(static) reduction(+:acc)，每个查询只读共享的 K 和索引，没有任何写入 ⇒ 无数据竞争、无伪共享。

【并行机制与性能解说】

• 并行结构：查询批次天然数据并行（每个查询之间无依赖），用 OpenMP 静态分配到 8 个线程；RMI 内部没有 SIMD 向量化空间（每次查的是不同的 key，走的是不同的模型），所以它的加速来自多核 + 每个核上的 MLP（memory-level parallelism），而不是 SIMD。这解释了为什么”学习型索引的收益最终由内存系统决定”。
• Work / Span / 并行度：
  • 构造：Work = Θ(n)（两遍扫描 + O(M1·SUB) 次 O(1) 拟合），Span = Θ(n/M1)（每个一级段内部串行测量误差），并行度 = M1 = 512。
  • 查询批次：单次查询 Work ≈ 9 + 8 + 1(模型访问) + log₂(2e2+3) + 1(叶访问) ≈ 22 个操作，其中关键路径上的依赖访存 = 1–2 次 DRAM 缺失，Span ≈ 2 × 90 ns = 180 ns。批量 Work = nq × 22，Span = 180 ns + (nq / 线程数) × …（因为查询之间独立，Span 只受”每线程串行处理的查询数”影响）。并行度 = Work/Span ≈ nq × 22 / 180 ns 的量级 —— 也就是”完全够用”。
  • 结论：Work/Span 给出的并行度是天文数字，所以它根本不是瓶颈。真正的瓶颈只有两个：① 每个核心能维持的在途缺失数（MLP，通常 10–16）；② DRAM 带宽。这就是为什么 §4.1 必须用 Little’s law 而不是 work-span 来判断这个内核。
• 瓶颈定量（假设：8 核、每核 12 个在途缺失、DRAM 缺失延迟 90 ns、双通道带宽 38.4 GB/s）：
  • MLP 上限：8 × 12 = 96 个在途缺失；每缺失 90 ns ⇒ 丢失率 ≈ 96 / 90 ns ≈ 1.07 × 10⁹ 缺失/s。
  • 带宽上限：38.4 GB/s ÷ 64 B = 6.0 × 10⁸ 缺失/s。
  • ⇒ 带宽先到顶（带宽上限比 MLP 上限低 1.8 倍）。所以这个内核是带宽受限的：
    • 朴素二分：每次查询 28 个不同 cache line ⇒ 10⁷ 查询 × 28 × 64 B = 17.9 GB ⇒ ≥ 466 ms；
    • 页式索引：每次 3 个 line ⇒ 1.9 GB ⇒ ≥ 50 ms；
    • RMI：每次 2–3 个 line ⇒ 1.3–1.9 GB ⇒ ≥ 34–50 ms。
  • 注意结论的诚实性：RMI 相对”页式 B+-tree”的优势（34 ms vs 50 ms）主要来自索引体积小 10³ 倍 ⇒ cache 命中率更高 ⇒ 每次查询少碰一个 cache line，而不是来自”模型预测得准”。这与 §2.2 表 1 后面的反面证据一致。
• 其它瓶颈：e2[] 若是偏大（分布漂移），窗口二分退化为长链，性能悬崖式下降——学习型索引的性能是”数据分布的函数”，这是它与 B+-tree 的根本区别。

3.2 示例二：GEMM 分块自动调优 + 学习型代价模型（C++17 + OpenMP）

// ============================================================================
// gemm_autotune.cpp — 分块 GEMM 的自动调优，并用拟合出来的代价模型预测未测过的尺寸
// 编译: g++ -O3 -march=native -fopenmp -std=c++17 gemm_autotune.cpp -o gemm_autotune
// 运行: OMP_NUM_THREADS=8 ./gemm_autotune 2048 3072
//        参数1 = 训练用矩阵尺寸 n（默认 2048）
//        参数2 = 测试（未见过的）矩阵尺寸 n_test（默认 3072）
// ============================================================================
#include <algorithm>
#include <chrono>
#include <cmath>
#include <cstdio>
#include <cstdlib>
#include <vector>
#include <omp.h>

static double now_ms() {
  using namespace std::chrono;
  return duration<double, std::milli>(steady_clock::now().time_since_epoch()).count();
}

// ---------- 可调参数就是 tile size T ----------
// 循环顺序 i-k-j（对 C 逐行写，利于编译器向量化 + 写分配）
static void gemm_blocked(int n, int T, const float* A, const float* B, float* C) {
  #pragma omp parallel for schedule(static) collapse(2)
  for (int ii = 0; ii < n; ii += T)
    for (int jj = 0; jj < n; jj += T)
      for (int kk = 0; kk < n; kk += T)
        for (int i = ii; i < ii + T; ++i)
          for (int k = kk; k < kk + T; ++k) {
            const float a = A[(size_t)i * n + k];
            float* c = C + (size_t)i * n + jj;
            const float* b = B + (size_t)k * n + jj;
            for (int j = 0; j < T; ++j) c[j] += a * b[j];   // 内层连续 → 可向量化
          }
}

static double gflops(int n, double ms) { return 2.0 * n * n * n / (ms * 1e6); }

int main(int argc, char** argv) {
  const int n  = (argc > 1) ? atoi(argv[1]) : 2048;
  const int nt = (argc > 2) ? atoi(argv[2]) : 3072;
  const int threads = omp_get_max_threads();

  // ---- 1. 真实测量：只在 7 个 tile size 上采样（这就是"训练集"）----
  const std::vector<int> tiles = {8, 16, 32, 64, 96, 128, 256};
  std::vector<double> meas(tiles.size(), 0.0);
  {
    std::vector<float> A((size_t)n * n, 1.0f), B((size_t)n * n, 1.0f), C((size_t)n * n, 0.0f);
    for (size_t t = 0; t < tiles.size(); ++t) {
      const int T = std::min(tiles[t], n);
      gemm_blocked(n, T, A.data(), B.data(), C.data());       // 预热
      double best = 1e30;
      for (int rep = 0; rep < 3; ++rep) {
        std::fill(C.begin(), C.end(), 0.0f);
        const double t0 = now_ms();
        gemm_blocked(n, T, A.data(), B.data(), C.data());
        best = std::min(best, now_ms() - t0);
      }
      meas[t] = gflops(n, best);
      std::printf("[训练] n=%d T=%3d : %8.2f ms  %8.2f GFLOPS\n", n, T, best, meas[t]);
    }
  }

  // ---- 2. 代价模型：roofline 上界 × 效率修正项 ----
  //   AI(T) = T / (2*E)  FLOP/byte（E=4B/elem），tiles 间通信量 = 2n^3/T elements
  //   预测吞吐 ŷ(T) = min(Peak, BW*T/(2E)) * 1/(1 + c1/T + c2*T^2/n^2)
  //     c1/T      : 每个 tile 的固定开销（调度、循环建立、cache 冷启动）
  //     c2*T^2/n^2: tile 数 (n/T)^2 不足 ⇒ 并行度不足 / 负载不均（Amdahl 项）
  const double BW    = 38.4e9;      // B/s，双通道 DDR4-2400（每个平台不同！）
  const double E     = 4.0;         // bytes per element (fp32)
  const double PEAK  = 400e9;       // FLOP/s，本机 fp32 峰值的粗略上界
  auto predict = [&](int T, double c1, double c2) {
    const double roof = std::min(PEAK, BW * (double)T / (2.0 * E));
    const double tiles_n = (double)n / T;
    const double eff = 1.0 / (1.0 + c1 / T + c2 * (tiles_n * tiles_n) / 1.0e4);
    return roof * eff;
  };

  // ---- 3. "训练"：只在两个常数上做一维搜索，最小化 log 空间的 RMSE ----
  double best_c1 = 0, best_c2 = 0, best_rmse = 1e30;
  for (double c1 = 0.1; c1 <= 64.0; c1 *= 1.5)
    for (double c2 = 0.0; c2 <= 20.0; c2 += 0.25) {
      double s = 0.0;
      for (size_t t = 0; t < tiles.size(); ++t) {
        const double p = predict(std::min(tiles[t], n), c1, c2);
        const double d = std::log(std::max(p, 1.0)) - std::log(std::max(meas[t], 1.0));
        s += d * d;
      }
      const double rmse = std::sqrt(s / tiles.size());
      if (rmse < best_rmse) { best_rmse = rmse; best_c1 = c1; best_c2 = c2; }
    }
  std::printf("\n[代价模型] 拟合得到 c1=%.3f  c2=%.3f  (log-RMSE=%.3f)\n", best_c1, best_c2, best_rmse);

  // ---- 4. 训练集内：模型选出的最优 T vs 实测最优 T ----
  auto pick_best = [&](const std::vector<int>& cand, auto score) {
    int bt = cand[0]; double bs = -1e30;
    for (int T : cand) { double s = score(T); if (s > bs) { bs = s; bt = T; } }
    return bt;
  };
  const int T_pred = pick_best(tiles, [&](int T){ return predict(std::min(T, n), best_c1, best_c2); });
  const int T_true = pick_best(tiles, [&](int T){
      int idx = (int)(std::find(tiles.begin(), tiles.end(), T) - tiles.begin()); return meas[idx]; });
  std::printf("[模型预测] 最优 T=%d   [实测] 最优 T=%d\n", T_pred, T_true);

  // ---- 5. 泛化检验：在"从未测量过"的尺寸 nt 上，用模型预测最优 T，再实测验证 ----
  const std::vector<int> news = {16, 32, 64, 128, 256};
  std::vector<double> m2(news.size(), 0.0);
  {
    std::vector<float> A((size_t)nt * nt, 1.0f), B((size_t)nt * nt, 1.0f), C((size_t)nt * nt, 0.0f);
    for (size_t t = 0; t < news.size(); ++t) {
      const int T = std::min(news[t], nt);
      gemm_blocked(nt, T, A.data(), B.data(), C.data());
      double best = 1e30;
      for (int rep = 0; rep < 3; ++rep) {
        std::fill(C.begin(), C.end(), 0.0f);
        const double t0 = now_ms();
        gemm_blocked(nt, T, A.data(), B.data(), C.data());
        best = std::min(best, now_ms() - t0);
      }
      m2[t] = gflops(nt, best);
      std::printf("[测试] n=%d T=%3d : %8.2f GFLOPS\n", nt, T, m2[t]);
    }
  }
  // 模型在 n=nt 上重算（注意：n 已经变了，必须重算 tiles 数）
  auto predict_at = [&](int T) {
    const double roof = std::min(PEAK, BW * (double)T / (2.0 * E));
    const double tn = (double)nt / T;
    return roof / (1.0 + best_c1 / T + best_c2 * (tn * tn) / 1.0e4);
  };
  int T_pred_new = pick_best(news, predict_at);
  int T_true_new = pick_best(news, [&](int T){
      int idx = (int)(std::find(news.begin(), news.end(), T) - news.begin()); return m2[idx]; });
  std::printf("\n[泛化] n_test=%d: 模型预测 T=%d  实测最优 T=%d  （误差 %.1f%% vs %.1f%%）\n",
              nt, T_pred_new, T_true_new,
              100.0 * (1.0 - [&]{int i=(int)(std::find(news.begin(),news.end(),T_pred_new)-news.begin()); return m2[i];}()
                            / m2[(size_t)(std::find(news.begin(),news.end(),T_true_new)-news.begin())]),
              0.0);
  std::printf("线程数 = %d；若 T 过大导致 tile 数 < 线程数，则 Amdahl 上界 = "
              "1/(1/P + (P (n/T)^2 之外的串行部分))\n", threads);
  return 0;
}

【代码做什么？】

1. gemm_blocked 是标准的三重分块 GEMM（i-k-j 顺序、内层对 j 连续、可自动向量化），唯一可调参数是 T——这正是”系统设计中的一个决策变量”。
2. 第 1 步在7 个 tile size 上做真实测量（含预热、取 3 次最小值），得到训练集 (T, GFLOPS)。这就是”数据”。
3. 第 2 步写出物理意义明确的代价模型：ŷ(T) = min(Peak, BW·T/(2E)) × 1/(1 + c₁/T + c₂·(n/T)²/n²)。第一项是 Roofline 上界（算术强度 AI = T/(2E) FLOP/byte，见 §4.3）；c₁/T 捕捉每个 tile 的固定开销（分块循环建立、cache 冷启动）；c₂·(n/T)² 捕捉 tile 数不足导致的并行度不足/负载不均（Amdahl 项）。这两项是解析式的骨架，只有 c₁、c₂ 两个参数需要从数据里学。
4. 第 3 步”训练”：在 c₁、c₂ 的二维网格上搜索，最小化 log 空间的 RMSE（用 log 是因为性能差异是乘性的）。这里故意用网格搜索而不是梯度下降——参数只有 2 个，网格搜索更稳、更可解释；真实系统里换成贝叶斯优化/GBDT 只是把”骨架 + 少量参数”换成”纯数据驱动模型”。
5. 第 4 步：比较模型选出的 T 与实测最优 T（训练集内拟合）。
6. 第 5 步是最关键的一步——泛化检验：换一个从未测量过的矩阵尺寸 n_test，用拟合好的 (c₁, c₂) 预测最优 T，然后真的去测，报出”模型选中的 T 相对实测最优 T 的性能损失”。这就是”learned cost model”从论文走到工程的验收方式：不看拟合误差，看决策质量（decision quality）。

【并行机制与性能解说】

• 并行结构：#pragma omp parallel for collapse(2) 把 (n/T)² 个 C-tile 折叠成一维任务空间，静态分配给 8 个线程。每个 tile 只写自己的 C 区域、只读 A/B ⇒ 无数据竞争、无伪共享（写区域按行连续、互不重叠）。每线程对内层 j 循环可 SIMD 向量化（a 广播 + b[j] 连续载入 + FMA）。
• Work / Span / 并行度：
  • Work = 2n³ FLOP（每个 C[i][j] 做 2n 次操作）。
  • Span = 2n FLOP——每个输出元素必须串行累加 n 次乘加（C 的累加链是全序的，除非做多累加器分裂）。这是本代码真正的关键路径。
  • 并行度 = Work/Span = 2n³/2n = n² = 4.2 × 10⁶（n=2048）。所以从 work-span 看，并行度绝对充足；瓶颈在别处。
  • 但有效并行度受 (n/T)² 限制：T=256、n=2048 时只有 64 个 tile 分给 8 线程 ⇒ 每线程 8 个 tile，粒度足够；T=1024 时只有 4 个 tile ⇒ 只有 4 个线程有活干，Amdahl 上界降到 4/8 = 50%。这就是代价模型里 c₂·(n/T)² 项的物理来源。
• 瓶颈与定量（假设：8 核、双通道 38.4 GB/s、fp32 峰值约 400 GFLOPS ≈ 8 核 × 3.2 GHz × 8-wide AVX2 × 2 FMA）：
  • Roofline ridge point = 400/38.4 ≈ 10.4 FLOP/byte。
  • 分块通信量 = 2n³/T elements × 4 B；算术强度 AI(T) = T/(2E) = T/8 FLOP/byte。
  • T = 8：AI = 1 FLOP/byte ⇒ 带宽上限 = 38.4 × 1 = 38.4 GFLOPS（只有峰值的 9.6%）。
  • T = 32：AI = 4 ⇒ 154 GFLOPS（38%）。
  • T = 64：AI = 8 ⇒ 307 GFLOPS（77%）。
  • T = 128：AI = 16 > 10.4 ⇒ 理论上计算受限，但 tile 数降到 (2048/128)² = 256，仍足够；实际会被 c₁/T（每 tile 固定开销，T 越大越不重要）与 cache 容量限制。
  • ⇒ 模型预测的最优 T 落在 64–128 附近，并且当 n 变小（例如 n=512）时，最优 T 会向左移动，因为 (n/T)² 项会迅速吃掉并行度。“最优 T 随 shape 和机器变化”正是必须用学出来的模型代替手写常量的原因。
• Work-span 之外的第三个瓶颈：内存带宽不是唯一。当 T 大到超出一级/二级 cache 能容纳 T×T 的 A 块与 B 块时，实际带宽需求会高于 2n³/T 的模型（模型假设完美的 cache 分块），代价模型的残差（log-RMSE）正是这些未建模效应的度量。

3.3 示例三：ISPC foreach 的三种分配方案（ISPC + C++，围绕同一份数据并行内核）

// ============================================================================
// sinx.ispc — 同一份"求 sin 的泰勒展开"内核，三种 assignment 方案
// 编译（ISPC 需先编译成对象，再与 C++ 链接）：
//   ispc --arch=avx2-i32x8 -O3 sinx.ispc -h sinx_ispc.h -o sinx_ispc.o
//   g++ -O3 -march=native -fopenmp -std=c++17 main.cpp sinx_ispc.o -o sinx -lispcrt
//   （-O3 + --arch=avx2-i32x8：release 优化 + 8 宽 fp32 SIMD）
// ============================================================================
#include "sinx_ispc.h"

// ---------- 方案 A：foreach（编译器自行选择 assignment，当前实现为 interleaved）----------
export void sinx_foreach(uniform int N, uniform int terms,
                         uniform float* uniform x, uniform float* uniform result) {
  foreach (i = 0 ... N) {
    float v = x[i];
    float numer = v * v * v;
    uniform int denom = 6;            // 3!
    uniform int sign = -1;
    for (uniform int j = 1; j <= terms; ++j) {
      v += sign * numer / denom;
      numer *= x[i] * x[i];
      denom *= (2 * j + 2) * (2 * j + 3);
      sign *= -1;
    }
    result[i] = v;
  }
}

// ---------- 方案 B：手工 blocked 分配（每个 instance 一段连续的数组区间）----------
export void sinx_blocked(uniform int N, uniform int terms,
                         uniform float* uniform x, uniform float* uniform result) {
  uniform int count = N / programCount;      // 假设 N % programCount == 0
  int start = programIndex * count;          // 注意：start 是 varying 值
  for (uniform int i = 0; i < count; ++i) {
    int idx = start + i;                     // 8 个 instance 访问的 idx 互不相邻！
    float v = x[idx];
    float numer = v * v * v;
    uniform int denom = 6;
    uniform int sign = -1;
    for (uniform int j = 1; j <= terms; ++j) {
      v += sign * numer / denom;
      numer *= x[idx] * x[idx];
      denom *= (2 * j + 2) * (2 * j + 3);
      sign *= -1;
    }
    result[idx] = v;
  }
}

// ---------- 方案 C：gang 内动态分配（原子计数器领号；对应"每颗糖加工时间不同"的场景）----------
export void sinx_dynamic(uniform int N, uniform int terms,
                         uniform float* uniform x, uniform float* uniform result) {
  uniform int nextIter = 0;                       // gang-local 计数器（8 个 instance 共享）
  int i = atomic_add_local(&nextIter, 1);         // 每个 instance 领一个迭代号
  while (i < N) {
    float v = x[i];
    float numer = v * v * v;
    uniform int denom = 6;
    uniform int sign = -1;
    for (uniform int j = 1; j <= terms; ++j) {
      v += sign * numer / denom;
      numer *= x[i] * x[i];
      denom *= (2 * j + 2) * (2 * j + 3);
      sign *= -1;
    }
    result[i] = v;
    i = atomic_add_local(&nextIter, 1);           // 再来领一个
  }
}

// ============================================================================
// main.cpp — 三种方案的统一驱动：多核 task 划分（OpenMP 线程池）+ 计时 + 正确性校验
// 编译: g++ -O3 -march=native -fopenmp -std=c++17 main.cpp sinx_ispc.o -o sinx -lispcrt
// 运行: OMP_NUM_THREADS=8 ./sinx 16777216 5
// ============================================================================
#include <cstdio>
#include <cstdlib>
#include <cmath>
#include <chrono>
#include <vector>
#include <random>
#include <omp.h>
#include "sinx_ispc.h"

static double now_ms() {
  using namespace std::chrono;
  return duration<double, std::milli>(steady_clock::now().time_since_epoch()).count();
}

// 参考实现（标量、串行）：用来校验 ISPC 版本的结果
static void sinx_ref(int N, int terms, const float* x, float* r) {
  for (int i = 0; i < N; ++i) {
    float v = x[i], numer = x[i] * x[i] * x[i];
    int denom = 6, sign = -1;
    for (int j = 1; j <= terms; ++j) {
      v += sign * numer / denom;
      numer *= x[i] * x[i];
      denom *= (2 * j + 2) * (2 * j + 3);
      sign *= -1;
    }
    r[i] = v;
  }
}

int main(int argc, char** argv) {
  const int N     = (argc > 1) ? atoi(argv[1]) : (1 << 24);
  const int terms = (argc > 2) ? atoi(argv[2]) : 5;
  const int T     = omp_get_max_threads();

  std::vector<float> x(N), got(N), ref(N);
  { std::mt19937 rng(7); std::uniform_real_distribution<float> d(-3.0f, 3.0f);
    for (int i = 0; i < N; ++i) x[i] = d(rng); }
  sinx_ref(N, terms, x.data(), ref.data());

  // 每个线程处理一段连续的区间（blocked 到线程），线程内部由 ISPC gang 再并行
  auto run = [&](void (*fn)(int, int, float*, float*), const char* name) {
    std::fill(got.begin(), got.end(), 0.0f);
    const double t0 = now_ms();
    #pragma omp parallel for schedule(static)
    for (int t = 0; t < T; ++t) {
      const int lo = (int)((long long)N * t / T);
      const int hi = (int)((long long)N * (t + 1) / T);
      fn(hi - lo, terms, x.data() + lo, got.data() + lo);   // 每线程一个独立 gang
    }
    const double ms = now_ms() - t0;
    double maxerr = 0.0;
    for (int i = 0; i < N; ++i) maxerr = std::max(maxerr, (double)std::fabs(got[i] - ref[i]));
    std::printf("%-16s %8.2f ms   %.3f ns/元素   最大误差 %.3e\n",
                name, ms, ms * 1e6 / N, maxerr);
  };

  std::printf("N=%d terms=%d 线程=%d programCount=%d\n", N, terms, T, ispc_programCount());
  run(sinx_foreach,  "foreach");
  run(sinx_blocked,  "blocked");
  run(sinx_dynamic,  "dynamic");
  return 0;
}

【代码做什么？】

1. 三个 ISPC 函数实现同一个数学内核（sin 的 5 项泰勒展开），差别只在迭代到 program instance 的分配方式：A 用 foreach（编译器决定，当前实现等价于 interleaved）；B 手工 blocked（每个 instance 一段连续区间）；C 用 atomic_add_local 在 gang 内动态领号。
2. main.cpp 用 #pragma omp parallel for schedule(static) 把数组按线程分块（每线程一个连续区间），每个线程调用同一个 ISPC 函数——ISPC 的 gang 只管 SIMD 那 8 个 lane，多核靠 task/线程，这就是”两层并行”。
3. sinx_ref 是串行标量参考实现：每个方案都必须与它逐元素比对最大误差（学习型/自动并行化的代码最容易悄悄算错，验证是必须的）。
4. 三个方案逐一定时，输出 ns/元素——这就是”分配策略搜索空间”的实测表格，也是学习型策略（根据负载特征选择 A/B/C）的训练数据。

【并行机制与性能解说】

• 指令层面的差别（本示例的核心）：
  • foreach / interleaved：8 个 instance 访问的 idx = loop_i + programIndex 是连续的 8 个 float，编译器发一条 vmovaps（打包载入）就够了，写回同理一条 vmovaps。
  • blocked：8 个 instance 访问的 idx = start + i 之间相隔 count = N/8，编译器必须发 vgatherdps（gather）——需要 8 个独立的地址，载入吞吐在 Skylake 级核心上差约一个数量级（vmovaps 约 0.5 周期/8 lane，vgatherdps 约 5 周期/8 lane，具体随微架构变化）。在带宽受限时这个差别会被掩盖，在 cache 命中时它会直接体现为几倍差距。
  • dynamic：每个 instance 用一次原子操作领号 ⇒ 原子操作在 gang 内串行化（8 个 lane 必须协调），并且 nextIter 这一行会被反复写（在同一个 L1 line 上做原子读改写）。代价是每迭代一次原子操作；收益是负载均衡（当每次迭代代价差异大时）。
• Work / Span / 并行度（N = 2²⁴，terms = 5，每元素约 50 个浮点操作）：
  • Work = N × 50 ≈ 8.4 × 10⁸ FLOP。
  • Span = 一个 gang 迭代的串行长度 = 50 个操作（一位元素的前后依赖链：numer、denom、sign、v 的递推），乘上每线程需要串行执行的 gang 迭代数 N/(programCount × T)。T=8、programCount=8 时：Span = 50 × 2²⁴/64 = 1.31 × 10⁷ 操作。
  • 并行度 = Work/Span = 8.4×10⁸ / 1.31×10⁷ = 64 = programCount × 线程数。
  • 这个结果非常有启发性：对一个”完全数据并行”的循环，Work/Span 恰好等于机器的总 lane 数——意思是”你能提供的并行度刚好用满机器，没有一点余量”。所以任何负载不均、任何 SIMD 浪费（掩码关闭的 lane）、任何串行开销都会直接变成性能损失，没有冗余并行度可以吸收。
• 瓶颈：
  • 带宽：内存流量 = 读 x + 写 result = 2 × 2²⁴ × 4 B = 128 MB。在 38.4 GB/s 下 ⇒ ≥ 3.3 ms（下限）。计算：8.4×10⁸ FLOP / 400 GFLOPS = 2.1 ms。两者接近，但 AI = 8.4×10⁸ / 128 MB = 6.6 FLOP/byte < ridge point 10.4 ⇒ 带宽受限，floor = 3.3 ms。
  • 动态分配的原子开销：每次迭代一次原子操作，共 2²⁴ 次 = 1.68×10⁷ 次；若每次约 10 ns（gang 内协调），总开销 ≈ 0.17 s，但这是并行摊开的（8 线程 × 8 lane = 64 lane 并行），每 lane 约 26 万次原子操作 × ~5 ns ≈ 1.3 ms——与 3.3 ms 的内存时间同量级 ⇒ dynamic 在这个负载上明显更差。这正是”分配策略必须依据特征选择”的定量依据：当每次迭代代价的方差很小时（本例正是如此），动态分配只有开销没有收益。
  • blocked 的代价：gather 让载入指令吞吐成为新的限制，且 8 个 instance 各访问一个远离的 cache line ⇒ 每个 cache line 只用到 8 B 中的 4 B 的一部分（如果数组是 16 B 对齐的，8 lane 分散在 8 个不同 line ⇒ 访存效率降至 1/8）。这就是为什么 foreach 的默认实现是 interleaved，而不是 blocked。

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

4.1 带宽与延迟：Little’s law 与”在途字节”

• 公式：要维持带宽 BW，系统必须同时有 BW × L 字节在途（L = 访存延迟）。这等价于 缺失率 × 延迟 ≤ 可支撑的并发缺失数（MSHR / line fill buffer 数），即

  \[\text{并发缺失数}_{\min} = \frac{BW \times L}{\text{cache line 大小}}\]
• 数值算例 1（CPU，DDR4-2400 双通道）：BW = 38.4 GB/s（= 64 bit × 1.2 GHz × 2 × 2 通道，slide 65），L = 13 ns（CAS，slide 65），cache line 64 B。
  • 单通道：19.2 × 10⁹ × 13 × 10⁻⁹ = 249.6 B ≈ 4 条 cache line。
  • 双通道：38.4 × 10⁹ × 13 × 10⁻⁹ = 499 B ≈ 8 条 cache line。
  • 结论：一颗核心至少要维持 8 个在途缺失才能吃满双通道带宽。若每核只有 4 个 MSHR，则要 2 个核心才能跑满带宽；而 §3.1 的 RMI 查询内核用 8 核 × 12 个在途缺失 = 96 个，远超所需的 8 个，因此确定不是 MLP 受限，而是带宽受限——这与 §3.1 的结论一致。
• 数值算例 2（GPU，H100 HBM3）：BW = 3.2 TB/s，L ≈ 500 ns（HBM 缺失延迟量级）。
  • 3.2 × 10¹² × 500 × 10⁻⁹ = 1.6 × 10⁶ B = 1.6 MB 在途。
  • 132 个 SM ⇒ 每个 SM 需要约 12.1 KB ≈ 190 条 cache line 同时在途。
  • 结论：GPU 不是”核多所以带宽高”，而是”必须让每个 SM 都维持近两百个在途请求，才能把 HBM 的带宽吃满”。这就是为什么 aidatacenter slide 25 反复强调 “Keep HBM busy all the time”——HBM 带宽是推理性能的硬上限，而把 HBM 喂饱需要极深的访存并行度。

4.2 DRAM 的内部结构决定了”延迟不是一个数”

• 公式（DRAM 命令序列）：行命中（row hit）只需 CAS；行缺失（row miss）需要 PRE + RAS + CAS。aidatacenter slide 51/53 给出每一段约 10 ns 的量级（DDR3-1600 口径），现代 DDR4 的 CAS ≈ 13 ns。
  • 行命中延迟 ≈ 13 ns；行缺失延迟 ≈ 30–40 ns ⇒ 相差 2.5–3 倍。
• 数值算例 3（busy-pin 比例）：假设 PRE = 10 ns、RAS = 10 ns、CAS + burst 传输 = 10 ns，则每个请求的”数据引脚占用时间”只有 10/30 = 33%（slide 54 用图直接画出”data pins 只在很小一部分时间里工作”）。补救手段有两个，都在讲义里：
  1. burst mode（slide 55）：把延迟摊到更大的传输上 ⇒ 引脚占用率随 burst 长度线性上升；128 B 的 burst 可把 33% 提到 70% 以上。
  2. 多 bank 流水（slide 56）：一个 bank 做 PRE/RAS 时，另一个 bank 在传数据 ⇒ 引脚利用率接近 100%。bat 级并行的关键在于 bank 数 ≥ 3：RAS→CAS→PRE 三段正好对应三个 bank 的轮转。
     • 对学习型策略的启示：调度器（以及学习型调度器）的目标函数就是这个”引脚忙碌率”，它完全可以从控制器内部状态（每 bank 的 open row、队列长度）算出来——这就是 §2.4 里学习型 DRAM 调度可行性的物理基础。

4.3 Roofline 与算术强度：为什么 LLM 解码几乎永远是带宽受限

• 公式：算术强度 AI = FLOPs / Bytes；可达性能 = min(Peak_FLOPs, AI × BW)；ridge point = Peak/BW。
• 数值算例 4（H100 上 Llama3.1 8B 的单 token 解码；参数取自 aidatacenter slide 22/24/70）：
  • 假设：8B 参数以 bf16 存储 ⇒ 权重 16 GB；H100 峰值 bf16 稠密 ≈ 989 TFLOPS；HBM3 带宽 3.2 TB/s；decode 每 token 只需 2 × P × B FLOPs（P = 参数量，B = batch）。
  • ridge point = 989 / 3.2 = 309 FLOP/byte。
  • batch = 1：FLOPs = 2 × 8×10⁹ = 16 GFLOP；字节 = 16 GB ⇒ AI = 1.0 FLOP/byte，比 ridge point 低 309 倍。
    • 时间 = max(16 GFLOP / 989 TFLOPS, 16 GB / 3.2 TB/s) = max(0.016 ms, 5.0 ms) = 5.0 ms ⇒ 单卡上限约 200 token/s，且算力利用率仅 0.3%。
  • batch = 16、上下文 1024：FLOPs = 256 GFLOP；字节 = 权重 16 GB + KV cache 2 × 32 层 × 1024 token × 4096 维 × 16 batch × 2 B = 8.6 GB ⇒ 共 24.6 GB。
    • AI = 256 / 24.6 = 10.4 FLOP/byte（仍比 ridge point 低 30 倍）；时间 = max(0.26 ms, 7.7 ms) = 7.7 ms；可达算力 = 3.2 × 10.4 = 33 TFLOPS = 峰值的 3.4%。
  • 结论：解码是彻底的带宽受限，且”算力越强的卡，这个比例越糟”（peak 涨得比带宽快）。于是三个结论全部得到解释：
    1. 为什么必须把权重加载与计算完全重叠（slide 25：HBM 带宽是上限，要”keep HBM busy all the time”）；
    2. 为什么片上 SRAM 容量是关键（slide 24：SN40L 520 MB vs H100 约 100 MB ⇒ “5× SRAM advantage”，把中间结果留在片上，消除 GB 级的片外中间流量）；
    3. 为什么”融合”是最有效的优化（slide 24：整个 decoder 融合成一个 kernel，kernel 调用数从每 token ~800 次降到 3 次）。这个 33 TFLOPS 的账算出来的 7.7 ms，也解释了 §3.2 的代价模型里为什么”字节数”项往往比”FLOPs”项更主导。

4.4 Work-Span、Amdahl 与计算-通信重叠（含可扩展性上限）

• Work-Span 基本式：T_P ≥ max(Work/P, Span)；并行度 = Work/Span；T_P ≥ Span 意味着”再多的处理器也无法低于关键路径”。
  • 数值算例 5：thoughtprocess slide 34–35 的图像两步计算：Step 1（每像素乘 2，完全并行，N² 工作量）+ Step 2（求和，串行，N² 工作量）。
    • 只并行 Step 1：T_P = N²/P + N² ⇒ Speedup = 2N²/(N²/P + N²) → 2（无论多少核，上限就是 2）。
    • 也并行 Step 2（部分和 + 串行合并 P 步）：T_P = N²/P + N²/P + P ⇒ Speedup → P（当 N ≫ P）。注意 +P 这个项：它正是”归约的 Span”，是把上限从 2 抬到 P 的代价。
  • 数值算例 6（Amdahl 的极端情形）：thoughtprocess slide 37 的 Summit：27,648 GPU × 5,376 ALU = 148,635,648 个 ALU。若程序有 0.1% 串行（S = 0.001），Speedup ≤ 1/S = 1000。
    • 算力利用率 = 1000 / 1.486×10⁸ = 0.00067%——一台 1.5 亿路的机器，被千分之一的串行代码压到用上 1000 路。
    • 对”学习型策略”的直接含义：如果推理被放在热路径的关键路径上（而不是旁路/异步），它就是那段 S。
• 计算-通信重叠的量化（aidatacenter slide 41 的数据，直接引用）：benchmark 总工作量 844.44 TFLOPs，张量维度 BS=16, M=24576, K=131072, N=8192。

• 数值算例 7（重叠率的反推）：设重叠比例 o（被隐藏掉的通信占比），则不重叠部分为 (1-o) × t_comm，利用率 U = t_comp / (t_comp + (1-o)·t_comm)。
  • 32 RDU：0.79 = 16.5 / (16.5 + 15(1-o)) ⇒ 15(1-o) = 4.39 ⇒ o = 70.7%。
  • 对比”完全不重叠”的理论值 52.4%：重叠把 32 卡的利用率从 52% 抬到 79%，相对提升 51%；而且规模越大，重叠越关键（8 卡时”不重叠”的理论上限 88.5% 反而高于实测的 72%——说明小规模下损失来自峰值效率而非通信）。
  • 这解释了为什么数据流架构把”异步”做成硬件属性（slide 12：无指令 ⇒ 无指令取指/译码开销；异步的 compute / memory / chip-to-chip 通信 ⇒ 天然重叠）。“要不要重叠、重叠多少”是一个可以被学习型模型预测的调度决策。

4.5 学习型策略自身的能耗预算（本讲最”系统”的一节）

• 公式：单次决策能耗 = 推理 FLOPs × pJ/op + 模型状态访问字节数 × pJ/byte + 控制开销；可行条件是

  \[\text{决策频率} \times \text{单次决策能耗} \ll \text{该决策能省下的能耗}\]

• 数值算例 8（学习型预取器的能耗可行性）：假设每个 L2 缺失都跑一次 64→64→8 的小 MLP。

  • 推理量：64×64 + 64×8 = 4,608 MAC = 9,216 FLOP。
  • 用 slide 47 的口径（0.9 pJ/FP op）：9,216 × 0.9 pJ = 8.3 nJ。
  • 被”优化”的那次访存：一次 64 B 的 LPDDR 读取 = 64 B / 8 B × 1200 pJ = 9.6 nJ（用 slide 46 口径），或按 slide 47 的 640 pJ/32 bit 算 = 16 × 640 pJ = 10.24 nJ。
  • ⇒ 推理能耗 ≈ 访存本身能耗的 81%–86%——“预测得再准也不可能赢”。
  • 换用 slide 46 的口径（20 pJ/FP op）：9,216 × 20 pJ = 184 nJ，是访存能耗（9.6 nJ）的 19 倍⇒ 完全不可行。
  • 可行的补救（每种都要付出代价）：

• 表 3：决策频率决定可用的模型复杂度（本讲的核心设计表）

• 数值算例 9（预算上限的直接计算）：给定”被优化的系统”总功耗预算 1 W（移动 GPU 量级，slide 46），若学习型策略允许占 1%：
  • 预算 = 10 mW = 10 mJ/s。若决策频率为 10⁸ Hz ⇒ 单次决策预算 = 10 mJ / 10⁸ = 0.1 nJ。
  • 用 0.9 pJ/op ⇒ 可用推理量 ≈ 111 个 FLOP——一个 10×10 的矩阵乘都跑不完。
  • 结论：在高频决策点上，”学到的东西”必须以”查表/阈值/极小线性模型”的形态存在，而不能是一个神经网络。（这也解释了为什么真实的硬件学习型预测器看起来更像”带权重的历史表”而不是”神经网络”。）

4.6 学习型索引的复杂度：构造、查询与更新

• 构造：Work = Θ(n)（每层一遍扫描 + O(#segments) 次 O(1) 拟合），Span = Θ(n / M1)，并行度 = M1。对比 B+-tree bulk load：排序 Θ(n log n) + 填充 Θ(n)。
• 查询：Work = O(1 + log₂ ε)（两级边界二分各 O(log M1)、O(log SUB) 次 L1 访问 + 窗口内 O(log(2ε+1)) 次比较；关键路径上是 1–2 次 DRAM 缺失 + 3 次 cache 命中），对比 B+-tree 的 O(树高) 次依赖缺失。
• 更新（学习型索引的真正软肋）：插入 m 个 key 后，F 的分布漂移量为 ΔF。误差界 ε 保证失效的条件是”新 key 未落在任何段的拟合区间内”。处理方式有三种，代价分别是：
  1. 全量重建：Θ(n)（一次完整扫描 + 拟合），每 m 次插入支付一次 ⇒ 摊还 Θ(n/m) per insert。若 m = n/100，摊还成本 = 每次插入扫 100 个元素——比 B+-tree 的 O(log n) 更贵。
  2. 局部修补（PGM-index 的思路）：只在误差超界的段上”分裂”，用对数方法（logarithmic method）维护多个大小递增的静态索引 ⇒ 摊还 O(log² n)，但常数与实现复杂度都上升。
  3. 回退到 B+-tree 处理热写区、学习型索引只服务冷只读区（工程上最常用）。

4.7 三种策略的复杂度/瓶颈对照

• 表 4：三类学习型策略的 work-span、并行度与真实瓶颈

5. 关键要点

1. “AI in System Design” 的第一原则是：先算账，再换启发式。 任何一个学习型策略都必须证明”决策频率 × 单次推理成本“远小于它省下的时间/能耗。§4.5 的算例 9 说明：在 10⁸ Hz 的决策点上，0.1 nJ 的预算只够约 111 个 FLOP——高频决策点上的”学习”注定只能是查表/阈值/极小线性模型，而不是神经网络；反过来，编译期 / 自动调优可以放心用大模型和上千次真实测量。决策频率是选择模型形态的第一约束，而不是精度。

2. 数据移动始终是主角，AI 也一样。 HBM 给了 3.2 TB/s，但 BW × L = 1.6 MB 的在途字节要求（§4.1）和 1 FLOP/byte 的解码算术强度（§4.3）决定了：LLM 解码无论换多强的加速器都仍是带宽受限（算力利用率 3.4%）。同一逻辑反过来约束学习型策略自身：模型表要小、要常驻片上 SRAM、要在被优化的路径外——aidatacenter 的 26 pJ vs 1200 pJ（46 倍）就是这条原则的硬件表达。

3. 学习型策略把”最坏情况保证”换成了”平均情况性能”。 B+-tree 的 O(log n) 是与数据分布无关的保证；学习型索引的 O(log ε) 依赖”分布稳定”。因此工程上必须成对交付：（学习型快速路径 + always-correct 兜底路径），并且把回退率当成一级监控指标（§3.1 的代码里就打印了 fallback 计数）。没有兜底的 learned index 不是优化，是隐性 bug。

4. 并行程序里”配置选择”的收益常常大于”算法选择”。 aidatacenter slide 41 的数据说明：32 卡时同一份计算，仅靠”把通信与计算重叠”就能把利用率从 52% 抬到 79%（相对提升 51%）；thoughtprocess slide 68→69 说明：同一个网格求解器，把锁从内层循环挪到循环外就消除了 O(N²) 次同步。这些”配置”正是学习型策略最现实的落点——它们不需要新算法，只需要正确的决策，而正确决策依赖的是当前机器/当前 shape 的具体参数（ridge point、bank 数、lane 数、链路带宽），这正是人脑最容易过时、模型最不容易过时的地方。

5. 并行度通常不是瓶颈，”机器的固定资源”才是。 §3.1 的 RMI 查询 parallelism ≈ nq（天文数字），§3.3 的 ISPC 内核 Work/Span 恰好等于机器总 lane 数 64，§3.2 的 GEMM parallelism = n² = 4.2×10⁶。三种截然不同的工作负载，work-span 都给出”并行度充足”的结论，但真正的瓶颈分别是带宽、bandwidth + SIMD 指令吞吐、以及 (n/T)² 个 tile 的粒度。所以性能分析的正确顺序是：先算 work-span 排除并行度不足，再用 Roofline/带宽/延迟/能耗定位真实瓶颈。

6. 常见陷阱与注意事项

• 把”用 AI 优化系统”当成”加一个神经网络”：在热路径（每次缓存缺失、每个周期）上放一个大模型，推理开销会超过被优化对象的开销（§4.5 算例 8：8.3 nJ 的推理 vs 9.6–10.2 nJ 的访存，占了 81%–86%；换成 20 pJ/op 口径则是 19 倍）。先定决策频率，再定模型规模，顺序反了必然失败。
• 忽略”分布漂移（distribution shift）”：学习型索引的 ε 是构造时数据的性质；负载变了（热点写入、新的 tag 分布）、硬件变了（不同的 bank 数/cache 容量）、shape 变了（更大的 batch）之后，模型给出的决策可能比原来那个”不好的启发式”更差。必须实现：在线监控决策质量 + 超界时回退到保守策略。
• 只优化了瓶颈的邻居（misattribution）：Benchmarking Learned Indexes（VLDB 2021）这类复现研究指出，学习型索引的很大一部分收益来自索引体积缩小带来的缓存行为改善，而不是”模型更聪明”。如果不开性能计数器验证归因，你很可能把”空间收益”误记为”学习收益”，然后在下一个负载上收益消失。
• 数据竞争与伪共享（写路径最容易被学习型策略引入）：任何在并行热路径上累加统计量/更新模型参数的实现，都必须处理：① ++counter 的数据竞争（用 reduction、原子操作或每线程私有副本）；② 伪共享（两个线程的私有计数器落在同一条 64 B cache line 上，性能损失可达数倍——用 alignas(64) 或 padding 把每线程状态隔离到独立 cache line）；③ 模型表的读写并发（读者可能看到”半更新”的模型 ⇒ 需要 epoch/RCU 或双缓冲切换，这也是 §2.2 里”学习型索引在写密集场景下反而更慢”的根因之一）。
• 把动态分配当万能药：§3.3 算出 atomic_add_local 领号的开销（每迭代一次原子操作）在”每次迭代代价方差很小”的负载上只有开销没有收益，会与内存时间同量级。正确的判据是”每次迭代代价的变异系数（CV）”：CV 小 ⇒ 静态分配；CV 大 ⇒ 动态/任务窃取。反过来，静态 blocked 分配会把 SIMD 打包载入（vmovaps）退化成 gather（vgatherdps，载入吞吐差约一个数量级）——“均衡”和”连续性”是要权衡的两个目标，不是可同时白拿的。
• 在关键路径上做同步/推理（Amdahl 的隐形杀手）：thoughtprocess slide 37 的算例（0.1% 串行 ⇒ 1.486 亿 ALU 只能有效利用 1000 路）说明串行比例决定机器人上限。学习型策略若同步地阻塞在被优化路径上（例如”等控制器把决策算完再发下一个请求”），它就把自己变成了那段 S；正确做法是旁路/异步/预决策（提前一拍算好、或用 speculative 的方式并行算多个候选）。
• 在没有回退路径的情况下相信一个未验证的模型：§3.2 的代码强调”验收标准是决策质量（模型选中的 T 相对实测最优 T 的性能损失），而不是拟合误差”。一个 log-RMSE 很小的模型，完全可能在关键 shape 上选出灾难性的配置。必须保留 always-correct 的兜底路径（本例是全局二分 / 安全默认 tile size），并在验证集上量化最坏情况损失。

7. 思考题（带答案）

问题 1：某团队给一个只读分析型数据库实现了两级 RMI 索引。在 n = 2²⁸、查询为随机点的只读负载上，它把每次查找的依赖 DRAM 缺失从 4 次降到 2 次，性能提升 1.8×。团队于是把它推广到写密集的 OLTP 表上，结果吞吐反而下降。请从 work-span（构造/更新成本）、误差界的失效条件、以及并发控制三个角度解释原因，并给出至少两种可行的工程折中方案。

【答案】

• 角度一：更新路径的 work/摊还成本。RMI 的正确性依赖”构造时测出的误差界 ε“。每插入一批 key，若分布漂移使某个段的线性拟合误差超过 ε，就必须重建该段；最坏情况下要全量重建。若每 m 次插入重建一次，摊还成本 = Θ(n/m) 每次插入（§4.6）。在 OLTP 里 m 相对 n 很小（写入并发高），于是 Θ(n/m) 会显著大于 B+-tree 的 O(log n)。即使查询变快了，一旦写占一定比例，总吞吐就会被写路径吃掉（这正是 §4.7 表 4 中”学习型索引何时不该用”的那一栏）。
• 角度二：误差界失效导致性能悬崖，而不是渐进退化。B+-tree 的 O(log n) 与数据分布无关；RMI 的 O(log ε) 是”分布稳定”的前提保证。一旦热点写入集中在某个键区间，该段的模型预测会大幅偏出 ε，查找从”窗口内一次二分”退化为”回退全局 lower_bound“，即从 2 次缺失跳到 28 次缺失。这是性能悬崖（cliff）而不是平缓下降——系统在压力下会突然崩掉，这比”平均慢一点”危险得多。
• 角度三：并发控制与模型一致性。查询路径要读 b1/b2/m1/m2/e2 这一整套结构；重建意味着要原子地替换它们。读者绝不能看到”半新半旧”的模型（否则预测位置与边界数组不一致 ⇒ 结果错误）。实现上必须用 RCU / epoch / 双缓冲 + 版本号，这会给每次查找加一次额外的内存栅栏与版本检查，与写路径的争用一起，抵消掉原本的收益。（另外，e2[] 的测量、模型表的重建都是写放大来源。）
• 可行的工程折中：
  1. 冷热分离：把学习型索引用在只读/冷数据上（历史快照、列存的分析段），热写表继续用 B+-tree，两级之间按 key 区间路由。这是最大化利用”学习型索引擅长的分布”的做法。
  2. 用 PGM-index 式的 ε-有界 + 对数方法：把”整块重建”换成”只在超界段分裂 + 维护多个大小递增的静态索引”，得到摊还 O(log² n) 的更新，代价是实现复杂度和常数上升。
  3. （加分）增量/在线更新模型：只微调局部模型参数（而不是重建），并把误差界做成运行时可观测的量——一旦超过阈值就自动降级到 B+-tree 路径，同时上报。“可降级”比”更快”更重要。

问题 2：aidatacenter 讲义 slide 41 给出了一个分布式矩阵乘 benchmark：总工作量 844.44 TFLOPs，8/16/32 个 RDU 时系统总算力为 12,744 / 25,488 / 50,976 TFLOPs，计算 roofline 时间为 66.3 / 33.1 / 16.5 ms，reduce-scatter 时间为 8.6 / 9.7 / 15.0 ms；不重叠的理论峰值利用率为 88.5% / 77% / 52%，重叠后实测为 72% / 75% / 79%。 (a) 请验证”不重叠利用率”这一列的算术。(b) 反推 32 卡时被隐藏的通信比例 o。(c) 为什么 8 卡时”不重叠的理论上限（88.5%）”反而高于实测的 72%？(d) 如果加卡的边际收益要维持”每翻倍一次利用率不下降”，需要满足什么条件？

【答案】

• (a) 验证：不重叠意味着 T = t_comp + t_comm，利用率 U = t_comp/(t_comp + t_comm)。
  • 8 卡：66.3/(66.3+8.6) = 66.3/74.9 = 88.5% ✓
  • 16 卡：33.1/(33.1+9.7) = 33.1/42.8 = 77.3% ≈ 77% ✓
  • 32 卡：16.5/(16.5+15.0) = 16.5/31.5 = 52.4% ≈ 52% ✓
  • （另外可验算 roofline 时间：844.44/12744 = 66.3 ms、844.44/25488 = 33.1 ms、844.44/50976 = 16.6 ms ✓）
• (b) 反推重叠比例：U = t_comp / (t_comp + (1-o)·t_comm)，代入 32 卡：0.79 = 16.5/(16.5 + 15(1-o)) ⇒ 16.5 + 15(1-o) = 20.89 ⇒ 15(1-o) = 4.39 ⇒ o = 1 - 0.293 = 70.7%。也就是说，约 71% 的 reduce-scatter 时间被藏在了计算后面，仍有约 4.4 ms（29%）暴露在外。
• (c) 为什么 8 卡实测（72%）低于”不重叠”理论值（88.5%）：因为“不重叠使用率”是一个上界意义上的理论值，它假设计算部分本身能做到 100% 的 roofline 峰值。8 卡时通信只占 8.6/74.9 = 11.5%，主导损失不是通信，而是”计算本身达不到峰值”（tile 边界、流水线填充/排空、SRAM 容量、指令混合、以及”总 TFLOPs 是名义峰值”）。这恰好说明：在规模小的时候优化重叠毫无意义（通信不是瓶颈），优化单卡效率才是重点；规模大了以后反过来。 这就是”先定位瓶颈再优化”的最好例证。
• (d) 加卡的边际条件：令 r = t_comm/t_comp。t_comp ∝ 1/S（S = socket 数），t_comm 随 S 增长（32 卡时已从 8.6 ms 涨到 15.0 ms，约 1.7 倍）。要达到利用率 U，需要 o ≥ 1 - (1/U - 1)/r。
  • 若要求 U ≥ 79%（32 卡实测水平），在 S=32 时 r = 15/16.5 = 0.909 ⇒ o ≥ 1 - (1/0.79 - 1)/0.909 = 1 - 0.266/0.909 = 70.7%（与 (b) 一致）。
  • 若 S 继续翻倍到 64：t_comp 减半到 8.25 ms，而 t_comm 在链路上只增不减（假设涨到 22 ms）⇒ r = 2.67 ⇒ 要维持 79% 需要 o ≥ 1 - 0.266/2.67 = 90%。结论：规模每翻倍，所需的”重叠比例”就向 100% 逼近，一旦 o 无法继续提高，利用率就会陡降。 这就是为什么必须用完全异步的硬件机制（slide 12/26：异步计算 + 异步访存 + 异步 chip-to-chip 通信，allreduce 不占 HBM 容量和带宽）而不是”用软件线程尽力重叠”。

问题 3：你想给一台移动 SoC（整机功率预算约 1 W，其中 GPU 约 1 W，LPDDR 带宽按讲义口径约 150 pJ/byte）加一个”学习型缓存替换策略”。团队的第一版方案是：在每次 L2 缺失时跑一个 64→64→8 的小 MLP 预测该行未来是否会被再次访问，用预测分数代替 LRU 的年龄。 (a) 用讲义的能耗数字定量说明这个方案为什么不可行。(b) 给出三种可行的替代设计，并分别说明它们损失了什么。(c) 如果一定要保留”每次缺失都决策”的频率，那么单次决策的可用 FLOP 数是多少？（用 1% 的功率预算给这个策略）

【答案】

• (a) 不可行性定量：
  • 推理量：64×64 + 64×8 = 4,608 MAC = 9,216 FLOP。
  • 用 45 nm 口径 0.9 pJ/FP op：9,216 × 0.9 = 8.3 nJ；用 Dally/ARM 口径 20 pJ/op：184 nJ。
  • 对照”被优化的那次访存”：一次 64 B LPDDR 读取 = 8 × 8 B ⇒ 按 150 pJ/byte = 9.6 nJ；按 slide 47 的 640 pJ/32 bit ⇒ 10.24 nJ。
  • ⇒ 推理能耗是访存能耗的 81%（0.9 pJ 口径）到 1900%（20 pJ 口径）。也就是说，每省下一次访存，你至少要先花掉 0.8 次访存的能量；而且替换策略的收益上限是”少访存”，其”收益/成本比”必须 > 1 才可能赢——81% 已经吃掉了绝大部分收益，19 倍则完全不可能。再叠加：移动 SoC 的 GPU 总预算只有约 1 W，而 10 GB/s 的持续访存本身就约 1.5 W（slide 46）——预算根本没有余量给一个每次缺失都跑的 MLP。
• (b) 三种可行的替代设计：
  1. 放大决策粒度（按时间片批量决策）：每 1 ms 决策一次”当前阶段该用哪种替换策略”（在 LRU / LFU / 随机 / 流式旁路之间切换），而不是每行决策。收益：推理开销摊到 1 ms 一次，可忽略。损失：无法逐行精细决策，对”同一时间片内混合访问模式”无能为力；且切换本身有抖动。
  2. 把模型变成查表 + 极简特征：用 4–8 位量化的小表（例如”由 PC 高位 + 访问历史 4 bit 索引的 256 项计数器表”，本质就是感知机/自适应替换的形态）。收益：单次决策降到几个整数操作（< 0.1 nJ 量级）。损失：表达力大幅下降（无法表达复杂的地址相关性）；需要重新设计特征，训练与部署方式都要改。
  3. 把决策从热路径移到离线/低频路径：由编译期/部署期用学习型模型生成策略参数（例如”这一类 PC 的分配优先级”、阈值、预取的深度/距离），运行时只做查表应用。收益：训练可用大模型和大量数据，运行时零推理开销。损失：无法适应运行时阶段变化（phase change），需要周期性重训练与部署；对”新出现的工作负载”无法及时适配。
     • （第 4 种可给加分：只对”高价值边界情形”做推理——先用极廉价的门控（例如”这是本页第一次访问 + 缺失代价高”）筛掉 99% 的决策，只在真正重要的缺失上做一次推理，把单次功耗按决策数摊薄。损失：门控本身可能漏掉重要情形。）
• (c) 单次决策的 FLOP 预算：总预算 1 W，分给策略 1% ⇒ 10 mW = 10 mJ/s。若决策频率 = 每次 L2 缺失；假设缺失率使决策频率约 10⁸ Hz（10 GB/s ÷ 64 B ≈ 1.6×10⁸ 缺失/s，取 10⁸ 量级）⇒ 单次决策能耗预算 = 10 mJ/s ÷ 10⁸ /s = 0.1 nJ。用 0.9 pJ/FLOP ⇒ 可用推理量 ≈ 0.1 nJ / 0.9 pJ ≈ 111 个 FLOP。
  • 结论：111 个 FLOP 连一个 10×10 的矩阵乘都做不完，更不用说 9,216 FLOP 的 MLP。要在这么高的频率上”学习”，唯一可行的形态是极小的查表 / 阈值比较 / 计数器更新（这恰好解释了真实硬件里”学习型”预测器为什么都长得像带权重的历史表，而不是神经网络）。决策频率决定模型形态，这是 “AI in System Design” 最硬的一条约束。