Add register allocation note series

13 files changed, 4365 insertions, 0 deletions

diff --git a/content/notes/register-allocation/_index.md b/content/notes/register-allocation/_index.md
new file mode 100644
index 0000000..fdd4ccb
--- /dev/null
+++ b/content/notes/register-allocation/_index.md
@@ -0,0 +1,7 @@
+---
+title: "寄存器分配"
+lang: zh-Hans
+url: "/notes/register-allocation/"
+description: "寄存器分配的基本概念、经典算法、现代机器模型，以及 CPU 和 GPU 工业实现。"
+license: CC-BY-SA-4.0
+---
diff --git a/content/notes/register-allocation/advanced-topics/index.md b/content/notes/register-allocation/advanced-topics/index.md
new file mode 100644
index 0000000..51b2c3d
--- /dev/null
+++ b/content/notes/register-allocation/advanced-topics/index.md
@@ -0,0 +1,443 @@
+---
+title: "第十一部分: 高级问题与协同优化"
+slug: "advanced-topics"
+lang: zh-Hans
+series: register-allocation
+weight: 120
+created: 2026-08-19
+updated: 2026-09-09
+license: CC-BY-SA-4.0
+---
+
+前面的章节已经把 register allocation 的主要机制拆开讨论过. Liveness 决定 value 在什么位置必须存在, interference 或 live interval 描述这些生命周期怎样竞争物理资源, allocator 再通过 coloring, priority assignment, eviction, splitting 和 spilling 完成实际分配. 到工业编译器中, 这些机制很少独立工作. Instruction scheduling 会改变 live range, instruction selection 会改变 register class 和 operand constraints, coalescing 会扩大 allocation object, spill rewrite 又会制造新的 temporaries. GPU 上还多出 occupancy 这一层反馈, register pressure 的变化可能直接改变 resident waves 或 warps 的数量.
+
+因此后端优化中经常遇到一种共同结构: 某个 pass 在自己的局部目标上做出了更好的选择, 却把问题转移给了另一个 pass. Scheduler 提前若干 independent instructions 可以增加 ILP, 同时可能提高 register pressure. Aggressive coalescing 可以删除 COPY, 也可能让两个原本较短的 live ranges 合并成一个难以分配的长 range. Instruction selector 选择一条静态指令数更少的特殊指令, 可能引入 fixed-register constraint, 让后面的 allocator 付出额外 spill cost. 高质量 code generation 因而依赖这些阶段之间交换足够的信息, 而不是简单地把每个阶段各自优化到极致.
+
+## 1 Instruction Scheduling 与 Register Pressure
+
+Scheduling 对 RA 最直接的影响来自 definition 和 use 的距离. 一个 definition 被提前以后, 如果它的 uses 没有同时提前, 对应 live range 就会延长. 一个 last use 被推迟也会产生相同效果. 当很多 values 同时被拉长时, maximum register pressure 就可能上升.
+
+考虑两种合法的 schedule:
+
+```text
+Schedule A:
+
+a = load A
+b = load B
+c = load C
+
+x = f(a)
+y = g(b)
+z = h(c)
+```
+
+以及:
+
+```text
+Schedule B:
+
+a = load A
+x = f(a)
+
+b = load B
+y = g(b)
+
+c = load C
+z = h(c)
+```
+
+第一种 schedule 可以让多个 memory operations 较早进入执行流水线. 对具有较长 memory latency 的机器来说, 这种安排可能增加 memory-level parallelism. 代价是 `a`, `b`, `c` 会在较长区域内同时 live:
+
+```text
+a: |-------------------|
+b:    |----------------|
+c:       |-------------|
+```
+
+第二种 schedule 把 definition 放得更靠近 use:
+
+```text
+a: |-----|
+b:        |-----|
+c:               |-----|
+```
+
+register pressure 明显下降, 但同时存在的 independent memory operations 也减少了. 所以 scheduler 面对的不是简单的 "尽量缩短 live range", 它还要考虑 latency, critical path, execution resources 和 ILP.
+
+可以从 scheduling point 上的 live set 观察这种变化. 假设一条 ready instruction $I$ 产生若干新 definitions, 同时消费一部分已经到达 last use 的 operands. 如果暂时忽略 register class 和 multi-register value, 局部 pressure 变化可以粗略理解成 $\Delta P\approx |Def(I)|-|Kill(I)|$. 定义越多, pressure 越容易上升; 一次结束多个旧 values 的生命周期, pressure 则可能迅速下降.
+
+实际 scheduler 不能只根据这个局部变化选择指令. 假设当前机器有 16 个可用 registers, pressure 从 8 增加到 9 可能没有任何实际后果, 从 16 增加到 17 却可能迫使 allocator 产生 spill. GPU 上还会出现更明显的离散阈值: pressure 从 64 增加到 65 可能让 register allocation 跨过硬件 granularity, 进而减少一个 resident wave. 因此 pressure-aware scheduling 更关心 peak pressure 以及它和资源阈值之间的距离.
+
+Scheduling 和 allocation 的先后顺序也带来不同信息. Pre-RA scheduler 面对 virtual registers, 没有受到具体 physical assignment 的限制, 所以拥有更大的 instruction-motion 空间. Post-RA scheduler 已经知道真实 physical registers 和机器 hazards, 但寄存器复用也会引入新的 dependencies. 两个本来独立的 virtual values 如果最终使用同一个 physical register 的不同时间段, post-RA scheduler 就不能任意打乱这种复用关系.
+
+这也是为什么工业后端通常同时存在 pre-RA 和 post-RA scheduling. 前者需要控制 pressure, 后者则在 allocation 已经确定以后处理更具体的机器执行问题.
+
+## 2 Splitting, Spilling 与 Rematerialization
+
+当 register pressure 超过可用资源时, allocator 可以改变 assignment, 也可以改变 allocation object 本身. Splitting, spilling 和 rematerialization 都属于后一类手段, 它们通过修改 value 在不同程序区域中的 location 来降低寄存器需求.
+
+假设一个 value 有很长的生命周期:
+
+```text
+v:
+|----------------------------------------------|
+
+              high pressure
+                 |--------|
+```
+
+真正困难的区域可能只有中间这一小段. 如果 allocator 把整个 `v` spill, definition 之后就写入 stack, 每个后续 use 再重新 load, 那么大量本来没有压力的区域也要承担 memory traffic. Splitting 可以把原来的 live range 切成多个 fragments:
+
+```text
+v0:
+|--------------|
+
+v1:
+               |----------|
+
+v2:
+                          |--------------------|
+```
+
+`v0` 和 `v2` 仍然可以驻留 register, `v1` 在高 pressure region 内使用另一个 physical register或进入 memory. 这样一次 allocation failure 被限制在真正发生资源竞争的区域.
+
+Split point 本身也有成本. 两个 fragments 如果 location 不同, 边界上需要 COPY, spill 或 reload. 因此 allocator 不只需要决定 "在哪里压力最大", 还要决定 transfer 放在哪里执行代价最低.
+
+Loop boundary 是典型例子. 假设一个 transfer 可以放在 loop body 中:
+
+```text
+loop:
+    ...
+    COPY
+    ...
+```
+
+也可以放在 loop entry:
+
+```text
+COPY
+loop:
+    ...
+```
+
+如果 loop 执行很多次, 两种方案的静态代码只差一个位置, 动态执行成本却可能相差几个数量级. 所以 global splitting 通常会结合 block frequency 和 loop structure, 尽量让稳定的 register assignment 覆盖热区域, 把 location change 推到较冷的边界.
+
+Spilling 的成本也和 target ISA 密切相关. 某些机器指令允许 memory operand, 那么一个 spilled value 的 reload 有机会直接 fold 到原指令中. 另一些 ISA 要求 operand 必须位于 register, 每次 use 前都需要独立 reload temporary. 同一个 abstract spill 在不同 target 上最终可能产生完全不同的 machine instruction sequence.
+
+Rematerialization 则利用 value 的可重新计算性. 假设:
+
+```text
+v = 42
+...
+use v
+```
+
+如果把 `v` 正常 spill, 可能产生:
+
+```text
+store 42, [slot]
+...
+t = load [slot]
+use t
+```
+
+而 immediate 很容易重新生成时, 可以直接得到:
+
+```text
+t = 42
+use t
+```
+
+这样 memory spill 完全消失. Frame address, constant pool address 和某些简单 arithmetic expression 也经常适合 rematerialization.
+
+Rematerialization 并非永远便宜. 如果重新计算 `v` 需要几个 source values, 那些 sources 必须在 rematerialization point 仍然可用, 它们的 live ranges 可能因此被延长. 一个看起来只有两条指令的 recomputation 也可能增加 dependency chain 或占用 execution ports. 所以 allocator 需要比较 reload cost 和 recomputation cost, 不能只判断 "这个 definition 能不能复制".
+
+Spill slot 本身还存在资源复用问题. 两个 spilled values 如果它们在 memory 中需要保存的时间不重叠, 可以共享同一个 stack slot:
+
+```text
+spill a:
+|----------|
+
+spill b:
+             |----------|
+```
+
+这种 stack-slot coloring 和 physical-register reuse 的基本思想相同, 只是资源从 register file 换成了 stack frame. Size, alignment 和 target addressing constraints 又会给它增加新的限制.
+
+## 3 Cost Model, Profile 与 Loop
+
+Register allocator 很少存在唯一的合法解. 大多数时候它面对的是一批都合法但成本不同的选择. 哪个 value 应该 spill, 哪个 physical register 更合适, 哪条 COPY 值得为了它扩大 live range, 某个 split boundary 应该放在哪个 block, 最终都需要 cost model.
+
+只统计静态 uses 很容易产生错误判断. 考虑:
+
+```text
+cold:
+    use a
+```
+
+和:
+
+```text
+hot_loop:
+    use b
+```
+
+两个 values 静态上各有一次 use, 动态执行频率却可能完全不同. 如果二者必须选择一个 spill, `b` 的一次 reload 可能在整个程序运行中执行数百万次.
+
+一个抽象的 spill cost 可以把每个 use 和 def 按执行频率加权. 例如可以把 reload 部分写成 $\sum_{u\in Uses(v)}Freq(u)\cdot Cost_{\mathrm{reload}}(u)$, 再加入 definitions 对应的 spill/store cost. 真正实现还会继续修正这个模型: 某个 use 可以 fold memory operand 时 reload cost 会下降, value 可以 rematerialize 时 memory cost可能被计算成本替代, live range跨 call时不同 registers 又有不同保存成本.
+
+没有 profile 时, loop depth 常被用作 execution frequency 的近似. 位于三层 loop 内的 instruction 通常比函数 entry 附近的一次性代码更热. 这种 heuristic 很有用, 但它始终只是结构上的预测. 一个深层 loop 可能位于极少进入的错误处理路径, 一个没有显式 loop 的 basic block 也可能因为函数被大量调用而极热.
+
+Profile-guided optimization 能够给 allocator 更接近实际执行的 block frequency 和 edge frequency. 这会直接改变 spill placement 和 splitting. 如果一个 live range 在热路径与冷路径发生冲突, allocator可以考虑保留热路径上的 register residency, 把 reload 或 COPY 推到冷 edge.
+
+Coalescing 同样可以使用频率信息. 一条位于函数初始化路径上的 COPY 即使保留下来也只执行一次, 为了删除它而合并两个大型 live ranges通常没有太大价值. 如果 COPY 位于一个频繁执行的 inner loop, register preference 就应该更强.
+
+Calling convention 也具有类似结构. 一个 value跨过多个 hot calls时, 把它放在 caller-saved register可能不断产生 save/restore. 使用一个 callee-saved register虽然会增加一次 prologue/epilogue 成本, 却可能降低整个函数中的动态 traffic. 如果这些 calls 都位于冷路径, 使用 caller-saved register又可能更划算.
+
+Cost model 无法精确预测现代 CPU 的所有行为. Out-of-order execution, cache, branch prediction, move elimination 和 memory-level parallelism 都会让一条静态 instruction 的实际成本随上下文变化. RA 阶段更实际的目标是获得可靠的相对排序. 它需要分辨明显昂贵和明显便宜的方案, 同时自身不能成为编译时间瓶颈.
+
+## 4 Instruction Selection 与 Register Allocation
+
+Instruction selection 决定了 allocator 最终面对什么样的机器指令. 两个完成相同语义的 instruction sequence 可以具有不同 register pressure, register classes 和 operand constraints, 所以 instruction selector 的选择会直接改变 RA 的问题形状.
+
+例如某个 operation 可以用普通三地址 sequence 完成:
+
+```text
+t1 = ...
+t2 = ...
+r  = OP t1, t2
+```
+
+也可以用一条特殊 instruction:
+
+```text
+r = SPECIAL_OP ...
+```
+
+后一种方案的静态 instruction count 更少, 但如果 `SPECIAL_OP` 要求固定 register pair, 或只能使用一个很小的 register class, allocator可能需要为了它移动其他 live values. 如果这个区域已经处于高 pressure 状态, 最终 spill code 的成本可能超过节省的一两条 instructions.
+
+Memory operand 也是 instruction selection 和 RA 联系很紧的地方. 某些 x86 指令允许一个 source 直接来自 memory. 一个将来发生 spill 的 value 如果能够直接成为 memory operand, allocator 可以省掉独立 reload. Selector 选择的 instruction form 如果不支持这种 encoding, 后面的 spiller 就失去 folding 机会.
+
+Two-address lowering 会更加直接地产生 register preference. 抽象三地址操作:
+
+```text
+v3 = ADD v1, v2
+```
+
+如果目标机器要求 destination 与第一个 source 相同, 后端可能转换成:
+
+```text
+v3 = COPY v1
+v3 = ADD v3, v2
+```
+
+当 `v1` 在 ADD 后死亡时, allocator很容易让 `v1` 和 `v3` 使用同一个 physical register, COPY 随之消失. 如果 `v1` 在后面仍然 live, 这个 COPY 可能必须真实存在. 因而 instruction selector 或 two-address lowering 选择哪一个 operand 与 destination tied, 可以根据 liveness 改变最终 copy cost.
+
+ABI 进一步为 selection 和 allocation 设置了函数边界约束. Arguments 和 return values 有固定 register convention, caller-saved 与 callee-saved registers具有不同生命周期成本. Tail call, varargs 和特殊 calling conventions 还会进一步限制可用 physical resources.
+
+从优化理论看, instruction selection, scheduling 和 register allocation 最好能够联合求解, 因为三个阶段会互相改变彼此的成本函数. 但把它们完全合并会使搜索空间迅速增长. 工业编译器通常仍然采用分阶段 pipeline, 再通过 register-pressure estimates, operand hints, rematerialization information 和 post-RA optimization 在阶段之间传递有限但高价值的信息.
+
+## 5 GPU 上的 Occupancy-aware 优化
+
+GPU 把 register pressure 的后果扩展到了并行执行资源. CPU 上一个函数多使用几个 registers通常主要影响它自己的 spilling 和 scheduling. GPU 上 per-thread 或 per-wave register usage 会决定一个 SM 或 CU 能同时驻留多少 execution contexts.
+
+因此 GPU 的 register-pressure cost 通常是非线性的. 假设某个硬件分配 granularity 使 64 个 VGPR 可以维持 4 个 resident waves, 65 个 VGPR 则只能维持 3 个. 把 pressure 从 70 降到 66 可能完全不改变 occupancy, 从 65 降到 64 却可能产生明显差异.
+
+这使 GPU scheduler 和 allocator 关心 occupancy threshold, 而不只关心 raw register count. 如果当前 pressure 已经远低于下一个 threshold, scheduler 可以更积极地提前 independent operations, 用较长 live ranges换取 ILP:
+
+```text
+load a
+load b
+load c
+
+compute a
+compute b
+compute c
+```
+
+当 pressure 接近 cliff 时, 同样的 scheduler 可能更倾向于把 definition 靠近 use:
+
+```text
+load a
+compute a
+
+load b
+compute b
+```
+
+这种策略降低 peak pressure, 但也减少同时进行的 independent work. 所以 GPU 上 ILP 和 occupancy 经常形成直接竞争.
+
+Spilling 也不能脱离 occupancy 评价. 一个 allocator 可以通过 spill 把 VGPR count 从 65 压到 64, 获得一个额外 resident wave. 如果只产生少量 scratch traffic, 最终吞吐可能提高. 如果为了这一点引入大量 scratch loads/stores, 更多 occupancy 也可能无法弥补 memory cost.
+
+AMDGPU 还存在多维 register pressure. SGPR, VGPR 和 AGPR 消耗不同资源, uniformity analysis 可以把某些 wave-uniform values 从 VGPR 转移到 SGPR. 当 VGPR 是 occupancy bottleneck 而 SGPR 仍然宽松时, 这种 scalarization非常有效. 如果 SGPR 自己已经接近限制, 继续转移只是把瓶颈换了位置.
+
+因此 GPU pressure 更适合看成一个向量, 例如 $(P_{\mathrm{SGPR}},P_{\mathrm{VGPR}},P_{\mathrm{AGPR}})$. Hardware resource model 再把这个向量映射成能够同时驻留的 waves 数量. Scheduler, register-bank selection 和 allocator 的决策都应该围绕这个整体资源状态展开.
+
+这也是 GPU RA 与 CPU RA 在高级优化阶段最大的差异之一. CPU 上 register count 更多表现为局部代码生成成本, GPU 上 maximum register usage 还能变成 kernel-level execution capacity.
+
+## 6 Machine Learning Guided Register Allocation
+
+Register allocation 长期依赖大量人工设计的 heuristics. Spill weight 决定哪些 live ranges 更值得保留, allocation priority 决定先处理谁, eviction heuristic 决定是否撤销已有 assignment, splitting heuristic 决定在哪里切开 live range, coalescing heuristic 则需要估计删除 COPY 与增加 register pressure 之间的收益. 这些决策都受到 CFG, block frequency, live-range shape, register class, call crossing 和 target architecture 的共同影响. 手工 heuristic 可以稳定运行, 但随着状态维度增加, 很难用一个简单公式完整描述 "在这种程序结构和机器上哪一个决策最终产生更快的代码". Register allocation 因而成为 machine learning 进入 compiler backend 的一个自然位置. LLVM 的 MLGO 就把 register allocation 的 Greedy eviction heuristic 作为 ML-guided optimization 的应用之一.
+
+RA 同时有一个很适合 compiler ML 的结构: allocation quality 可以交给 learned policy 预测, machine legality 却可以继续由传统编译器精确维护. 假设 allocator 在状态 $s$ 下存在一组候选动作 $A(s)$, compiler 可以先根据 register class, interference, subregister alias, fixed-register constraint 等信息计算合法集合 $A_{\mathrm{legal}}(s)$, 模型只在这个集合中选择 $a=\pi(s)$. 这样模型可以学习 "哪一个合法选择更好", 却没有权力产生一个明显违反 ISA constraint 的 assignment. RL4ReAl 对 register type, physical-register congruence 和 interference 都建立了显式约束, 并通过限制 agent 的 action space 保证 coloring action 属于合法 register 集合.
+
+```text
+allocator state
+      |
+      v
+feature / graph representation
+      |
+      v
+learned policy
+      |
+      v
+choose an action
+      |
+      v
+compiler legality constraints
+      |
+      v
+update allocation state
+      |
+      +--------------------> next decision
+```
+
+LLVM MLGO 对 Greedy allocator 的改造很好地体现了这种思路. 第九部分已经看到, Greedy 为一个 live interval 尝试 physical register 时, 如果该 register 已经被其他 virtual live ranges 占据, allocator 可以考虑 eviction. 手工 Greedy 需要根据 interfering intervals 的 spill weight, hints 和其他状态判断某个候选 register 是否值得争夺. MLGO 没有重新设计一套 register allocator, 而是在这个 decision point 引入 `RegAllocEvictionAdvisor`: Greedy 仍然负责 `LiveInterval`, `LiveRegMatrix`, candidate physical registers, splitting, spilling 和 allocation state 的维护, advisor 负责从可行 eviction choices 中作出策略选择. LLVM 的 ML eviction implementation 使用 `MLModelRunner` 接收特征, 模型输出的 decision 名为 `index_to_evict`, 然后 allocator继续执行相应的 eviction 流程.
+
+
+```text
+                    LLVM Greedy
+
+LiveInterval
+    |
+    v
+try physical register R
+    |
+    v
+find interference
+    |
+    v
++-------------------------------+
+| RegAllocEvictionAdvisor       |
+|                               |
+| allocator state -> ML model   |
+|                  -> decision  |
++-------------------------------+
+    |
+    v
+evict / reject candidate
+    |
+    v
+Greedy continues
+    |
+    +--> splitting
+    +--> recoloring
+    +--> spilling
+```
+
+这种做法的工程价值在于可以替换一个复杂 heuristic, 同时保留几十年积累下来的 allocator infrastructure. 如果模型选择了一次收益不佳的 eviction, 结果通常表现为代码质量下降; interference checking 和 Machine IR legality 并不会因为预测模型本身缺少形式化理解而一起消失. ML model 因而更像 allocator 中的 advisor, 而不是拥有全部 machine-state authority 的 code generator. LLVM CodeGen 对 learned RA decision 也采用这种 advisor abstraction.
+
+真正困难的是怎样表示 allocator state. 一个 eviction decision 的价值不只取决于当前 interval 和 victim 的长度. 长 live range 可能只在冷 block 中发生 conflict, spill weight 高的 interval 也可能拥有便宜的 rematerialization, 一个 victim 被 eviction 后还可能重新进入 queue 并驱逐其他 intervals. 模型需要从有限 feature vector 中捕获这些后续效应. LLVM 的 ML RA infrastructure 会向模型提供 live-range 和 interference 相关的数据, 并能够利用 MachineBlockFrequencyInfo, MachineLoopInfo 和 register-allocation state 等 CodeGen 信息. 这种 structured-feature 方法的优点是 inference 很轻, 也容易嵌入已有 allocator; 缺点是模型只能利用 compiler writer 预先选择出来的状态描述.
+
+RL4ReAl 选择了更大的学习边界. 它没有只学习 Greedy 的某一个 heuristic, 而是把 register allocation 建模成 hierarchical multi-agent reinforcement learning. LLVM Machine IR 被转换成 interference graph, graph 中的 machine-level entities 使用 MIR2Vec 表示, 然后多个 agents 协同完成 node selection, task selection, live-range splitting 和 coloring. Coloring agent 在没有合法 register 时也负责进入 spilling decision. Splitter 选择 live range 的 split point, compiler按照这个选择真正修改程序, 更新后的 interference graph 再返回给 policy继续决策.
+
+其控制过程可以近似理解成:
+
+```text
+interference graph
+        |
+        v
+   Node Selector
+        |
+        v
+   Task Selector
+      /      \
+     v        v
+ Coloring   Splitting
+    |           |
+    |           v
+    |      modify live range
+    |           |
+    +-----------+
+        |
+        v
+updated interference graph
+        |
+        v
+next RL decision
+```
+
+RL4ReAl 的 coloring action 并不是任意预测一个 register number. Compiler 根据 virtual register type, physical-register overlap 和 interference 计算合法 register 集合, coloring agent 从这个集合中选择 assignment; 如果集合为空则产生 spill. Splitting agent 则根据 value representation, 各 use 位置的 spill weight 和 use distances 选择 split point. 这种设计把第四部分的 coloring, 第六部分的 splitting 以及 spill selection放进了一个连续决策过程, 因而能够学习传统 allocator 中多个 heuristics 之间的相互作用. 论文在 x86-64 和 AArch64 上使用 SPEC CPU 2006/2017 进行实验, 并报告其方案能够达到或超过所比较的 LLVM production allocators.
+
+从 ML 的角度看, RA 比普通 classification 问题更接近 sequential decision making. 假设 allocator 在时间 $t$ 做出 assignment $a_t$, 这个选择会改变后面的 interference state $s_{t+1}$; 一次不理想的 eviction 可能直到几十次 allocation decisions 后才导致 spill. 最终程序运行时间更是所有 decisions 共同作用的结果. 因此单独给每个 decision 构造一个精确的 supervised-learning label 很困难. RL 可以通过 reward 学习长期效果, 但同时引入 credit assignment 和训练成本问题. RL4ReAl 采用 hierarchical agents 和局部 reward, 其中 coloring reward与 LLVM 计算的 spill weight 联系, splitting reward则利用 splitting 前后 spill-weight state 的变化, 再配合 global reward协调不同 agents.
+
+Machine learning 还可以把学习边界继续推向完整 assignment. VeriLocc 研究了利用 LLM 直接进行 GPU register allocation: 模型从 machine-level representation 生成 target-specific register assignments, static analysis负责跨架构表示与约束信息, 随后 verifier检查生成结果; 验证失败时重新生成 assignment. 这与 advisor 模式相比已经接近 end-to-end learned allocator, 但 correctness 仍然由独立 verifier兜底. 该工作的实验集中在 GPU GEMM 和 multi-head attention kernels, 论文报告 85%-99% 的 single-shot valid allocation rate, `pass@100` 接近 100%, 并在部分 case study 中得到超过 rocBLAS 10% 的运行时间改善. 这些结果说明 learned policy 可以探索人工 heuristic 不容易覆盖的 assignment space, 同时也说明 verifier 在生成式 allocator 中承担了不可缺少的角色.
+
+
+```text
+hand-written allocator
+        |
+        v
+learn one heuristic
+MLGO eviction advisor
+        |
+        v
+learn several interacting decisions
+RL4ReAl
+        |
+        v
+predict complete assignments
+VeriLocc-style generation
+```
+
+学习边界越大, 模型能够重新发现的策略越多, 需要解决的训练和工程问题也随之扩大. Eviction advisor 每次只处理 Greedy 已经构造好的局部 decision point, feature 和 action space都相对受控. 完整 learned allocator 必须面对 live-range splitting带来的动态图变化, machine constraints, architecture portability以及长序列 decision 的 reward attribution. 如果进一步直接生成完整 assignment, verifier成本和失败后的搜索策略又成为整个系统的一部分.
+
+RA 中的 ML 还受到一个非常现实的约束: inference 本身属于 compiler compile-time. 一个运行数毫秒的模型如果在一次函数 allocation 中调用数千次, 总成本很容易超过原来的 heuristic. 同时, training distribution 与真实 workload 之间可能存在差异, architecture变化也会改变 register classes, instruction constraints 和 spill costs. 因此高质量 learned allocator 不能只追求 benchmark 上的 reward, 还需要考虑模型尺寸, inference frequency, feature extraction cost, generalization以及 fallback behavior.
+
+从编译器架构的角度看, 最有价值的模式仍然是把精确分析和 learned decision结合起来. Liveness, interference, register class和machine legality继续由确定性算法计算, ML 处理那些具有巨大搜索空间并长期依赖人工 heuristic 的 optimization choices. 这种划分既可以像 MLGO 一样只替换 Greedy 的 eviction policy, 也可以像 RL4ReAl 一样扩大到 splitting 和 coloring, 再进一步发展成带 verifier 的完整 assignment generation. Register allocation 因此也是观察 "ML 怎样真正进入工业编译器" 的一个很好的案例: ML 并不需要替代整个 compiler backend, 它首先可以替代 backend 中最难手工调优的决策函数.
+
+## 7 怎样评价 Register Allocation
+
+一个 register allocator 不能只通过 "spill 了多少 virtual registers" 来评价. Spill 的动态位置, reload folding, COPY 数量, frame size, callee-saved register使用情况以及 scheduler interaction 都会影响最终代码. GPU 还需要把 register count, scratch traffic 和 occupancy一起看.
+
+Compile-time evaluation 首先需要测 allocator 自身的开销. 对大型函数来说, interference queries, splitting, recoloring 或 graph construction 都可能显著消耗时间和内存. 一个 allocator 在 benchmark 上得到少量 runtime improvement, 如果后端编译时间增加很多, 是否值得采用取决于编译器的使用场景. JIT, interactive build 和离线 HPC compiler 对这种 trade-off 的接受程度完全不同.
+
+Generated-code evaluation 则至少要同时观察静态和动态结果. Static metrics 可以包括 code size, COPY 数量, spill stores, reloads, stack frame size 和 callee-saved register使用数量. Runtime measurement 才能判断这些变化经过 processor pipeline 和 memory hierarchy 后是否真正有收益.
+
+GPU 还应记录 per-thread register usage, scratch/local memory usage 和 occupancy-related resource data. 一个新的 allocator把 VGPR 从 72 降到 68 如果没有跨过任何 occupancy threshold, 这项变化本身可能意义有限. 如果从 65 降到 64 恰好增加一个 resident wave, 即使 register count只减少一个, 影响也可能更大.
+
+Benchmark 需要覆盖不同的 allocation difficulty. Straight-line arithmetic主要测试 local assignment 和 machine constraints. 大型 CFG 更容易暴露 global liveness 和 splitting 问题. Loop-heavy code会放大 hot-path spill placement. Vectorized code可以制造宽 register class pressure. 调用密集函数则适合测试 caller/callee-saved decisions. GPU kernel还需要区分 compute-bound, memory-latency-bound 和 occupancy-sensitive workloads.
+
+微基准可以专门隔离某个机制. 例如构造许多同时 live 的 values 测试 spill heuristic:
+
+```text
+v1 = ...
+v2 = ...
+v3 = ...
+v4 = ...
+...
+use v1
+use v2
+use v3
+use v4
+```
+
+构造一条长 global live range与大量短 local ranges可以观察 allocation order和eviction:
+
+```text
+global:
+|----------------------------------------------|
+
+local1:    |---|
+local2:          |---|
+local3:                |---|
+local4:                      |---|
+```
+
+大量 COPY 可以测试 coalescing, fixed-register instructions 可以测试 machine constraints, loop boundary 上的冲突则可以测试 splitting是否把 transfer放到了合理位置.
+
+真实应用 benchmark 仍然不可替代. 微基准能够证明某个机制按照预期工作, 无法告诉我们多个 heuristics结合以后是否在完整程序中稳定获益. 一个 allocator可能在多数程序上变化很小, 却在少数 high-pressure workloads上产生大幅退化. 这些 outliers 往往最有分析价值, 因为它们能暴露 cost model, splitting strategy 或 machine constraint handling 中的系统性问题.
+
+因此评价 RA 更适合看一组相互关联的数据: compile time, code size, spills, copies, register pressure, runtime, 以及 GPU 上的 occupancy 和 scratch traffic. 最终目标仍然是机器代码的整体性能, allocator内部某一个统计数字只能解释这个结果的一部分.
+
+到这里, register allocation 的算法, 工业实现和跨阶段优化已经形成一套相对完整体系. 第十二部分将脱离概念综述, 从一个最小 Machine IR 开始真正实现一个 register allocator: 先建立 CFG 和 liveness, 再构造 live ranges 与 interference, 完成一个可以工作的基础 allocator, 加入 spill rewrite, 最后逐步扩展到 coalescing, register classes 和更真实的机器约束.
diff --git a/content/notes/register-allocation/coalescing/index.md b/content/notes/register-allocation/coalescing/index.md
new file mode 100644
index 0000000..58a0f97
--- /dev/null
+++ b/content/notes/register-allocation/coalescing/index.md
@@ -0,0 +1,355 @@
+---
+title: "第五部分: Coalescing 与 Iterated Register Coalescing"
+slug: "coalescing"
+lang: zh-Hans
+series: register-allocation
+weight: 60
+created: 2026-08-14
+updated: 2026-09-09
+license: CC-BY-SA-4.0
+---
+
+第四部分暂时把 interference graph 当成了一张只包含硬约束的图: 两个节点之间有边, 就必须分配不同颜色. 真实机器 IR 中还大量存在另一类关系, 即 COPY 或 move. 对于一条 `b = COPY a`, allocator 希望 $a$ 和 $b$ 最终获得同一个物理寄存器, 这样 COPY 就可以消失. 这使寄存器分配同时面对两种方向相反的约束: interference 希望两个节点分开, move 希望两个节点靠拢.
+
+COPY 的来源很多. SSA destruction 会把 phi 转换成边上的 copy, calling convention 会在参数寄存器和普通虚拟寄存器之间产生 copy, instruction selection 和 two-address lowering 也会引入 move. Live range splitting 本身还会制造新的 copy, 用于连接同一个逻辑值的不同 fragments. 因此 allocator 如果完全忽略 move, 最终机器代码里往往会保留大量本可消除的寄存器间复制.
+
+让 COPY 两端获得同一个颜色的操作称为 register coalescing. 问题在于, coalescing 会合并两个 live ranges, 也会合并它们各自的 interference relation. 一次看起来很有收益的 move elimination, 可能让 interference graph 变得更难着色. Iterated Register Coalescing, 简称 IRC, 就是在 simplify, coalesce, freeze 和 spill 之间反复切换, 尽量消除 move, 同时控制 coalescing 对 colorability 的破坏.
+
+## 1 为什么不能看到 COPY 就直接合并
+
+考虑:
+
+```text
+v = COPY u
+```
+
+如果 $u$ 和 $v$ 没有 interference edge, 最直观的想法就是把两个节点合成一个节点 $uv$. 合并之后所有对 $u$ 或 $v$ 的引用都视为对同一个节点的引用, 如果最终 $uv\mapsto R_1$, 原来的 COPY 就会变成 `R1 = COPY R1`, 随后删除.
+
+但 "没有 interference edge" 只说明 $u$ 和 $v$ 的生命周期允许共享寄存器, 并不能保证整张图仍然存在一个让它们共享颜色的 $K$-coloring.
+
+假设 $K=3$, 有三个节点 $a$, $b$, $c$ 两两冲突, 因而形成一个三角形. 此外 $u$ 和 $a,b$ 冲突, $v$ 和 $b,c$ 冲突, $u$ 与 $v$ 之间存在 COPY, 但没有 interference edge:
+
+```text
+       u           v
+      / \         / \
+     a---b-------?   c
+      \___________/
+
+a, b, c form a triangle
+
+u interferes with a, b
+v interferes with b, c
+u <---- COPY ----> v
+```
+
+更准确地列出边就是:
+
+```text
+a -- b
+b -- c
+a -- c
+
+u -- a
+u -- b
+
+v -- b
+v -- c
+```
+
+三角形 $a,b,c$ 在三个寄存器下必须使用三种不同颜色. 假设 $a\mapsto R_0$, $b\mapsto R_1$, $c\mapsto R_2$. 那么 $u$ 与 $a,b$ 冲突, 所以只能使用 $R_2$; $v$ 与 $b,c$ 冲突, 所以只能使用 $R_0$. 原图完全可以 3-color, 只是 $u$ 和 $v$ 必须使用不同颜色.
+
+如果强行把 $u$ 和 $v$ 合并成 $w$, 那么 $w$ 会同时与 $a$, $b$, $c$ 冲突. 合并后的 $a,b,c,w$ 构成 $K_4$, 需要四种颜色. 原来能够在三个寄存器中完成分配的图, 经过这次 coalescing 后就无法再 3-color.
+
+所以 coalescing 需要一个保守条件. allocator 希望确认这次合并不会明显破坏后续 coloring 的机会, 然后才真正把两个节点收缩成一个.
+
+这种策略通常称为 conservative coalescing.
+
+## 2 Briggs criterion 与 George criterion
+
+Briggs criterion 从合并之后节点的 high-degree neighbors 数量出发. 假设准备合并 $u$ 和 $v$, 目标机器有 $K$ 个颜色. 先取两个节点当前邻居集合的并集 $Adj(u)\cup Adj(v)$, 再观察其中有多少节点满足 $degree(t)\ge K$. 如果这样的 high-degree neighbors 少于 $K$ 个, Briggs 认为这次合并足够保守.
+
+也就是说, Briggs 检查的是 $Adj(u)\cup Adj(v)$ 中 significant nodes 的数量, 其中 significant 通常指 $degree\ge K$ 的节点. 条件可以在正文里写成: 如果 $\left|{t\in Adj(u)\cup Adj(v)\mid degree(t)\ge K}\right|<K$, 那么允许 coalesce.
+
+它背后的考虑和 simplify 一致. 合并后的节点即使当前 degree 很高, 其中许多邻居如果本身属于 low-degree nodes, 后续 simplify 时还可能陆续被删除. 真正容易把合并节点困住的是那些 $degree\ge K$ 的邻居. 如果这种邻居还不到 $K$ 个, 合并后的结构通常仍然保留足够好的可着色性.
+
+刚才那个失败的例子正好可以用 Briggs criterion 检查. 在 $K=3$ 时, $a,b,c$ 都至少有 3 个邻居, 所以它们都是 significant. 对 $u$ 和 $v$ 做邻居并集后, high-degree neighbors 至少包含 ${a,b,c}$, 数量正好是 3, 不满足 "少于 $K$" 的要求, 因此 Briggs 会拒绝这次 coalescing.
+
+George criterion 更适合处理一个普通 virtual register 与 pre-colored physical register 的 coalescing. 例如有:
+
+```text
+v = COPY RAX
+```
+
+如果可以把 $v$ 合并到 pre-colored 的 `RAX` 节点, COPY 就能删除. 但 `RAX` 的颜色已经固定, 不能像普通节点一样参与随后的 simplify 和重新着色.
+
+设 $u$ 是 pre-colored node, 准备将普通节点 $v$ 合并到 $u$. George 的思路是逐个检查 $v$ 的邻居 $t$. 如果 $t$ 本身是 low-degree, 即 $degree(t)<K$, 那么它以后比较容易处理; 如果 $t$ 已经和 $u$ interfere, 合并不会给 $u$ 增加一条新的约束; 对 pre-colored 邻居也可以按固定颜色约束处理. 只有那些 high-degree, 尚未和 $u$ 冲突, 并且会因为合并而成为 $u$ 新邻居的节点, 才真正增加风险.
+
+在经典 IRC 的写法中, 对每个 $t\in Adj(v)$ 检查类似条件: $degree(t)<K$, 或者 $t$ 已经是 pre-colored, 或者 $(t,u)\in E$. 如果所有邻居都通过检查, 就可以把 $v$ 合并进 $u$.
+
+Briggs 更像是从合并后的整体邻域判断风险, George 更关注这次合并究竟给一个已有节点增加了哪些新的困难约束. 在 IRC 中, 当一端是 pre-colored node 时通常使用 George-style test; 两端都是普通节点时通常使用 Briggs-style conservative test.
+
+## 3 Coalescing 会怎样修改图
+
+假设经过 conservative test 后决定把 $v$ 合并到 $u$. allocator 不能只记录 "以后让它们用同一个颜色", 因为从这一刻开始, 其他图操作需要把它们看成同一个逻辑节点.
+
+如果原来:
+
+```text
+Adj(u) = {a, b}
+Adj(v) = {b, c}
+```
+
+合并以后大致得到:
+
+```text
+Adj(uv) = {a, b, c}
+```
+
+同时所有原来指向 $v$ 的 interference relation 都要转移到代表节点 $u$ 上. $u$ 和 $v$ 各自关联的 move 也要合并, 因为新的联合节点仍然可能与其他节点存在 COPY relation.
+
+经典实现通常不会真的把所有数据结构里的 `v` 文本替换成 `u`. 它会保留一个 `alias` 关系. 如果 $v$ 被合并进 $u$, 可以记录 $alias[v]=u$. 后面如果又把其他节点合并到 $v$, 查询时通过 `GetAlias(v)` 一直沿 alias 链找到当前代表节点.
+
+例如:
+
+```text
+v3 -> v2
+v2 -> v1
+```
+
+那么:
+
+```text
+GetAlias(v3) = v1
+```
+
+最终 coloring 完成以后, coalesced nodes 直接继承其代表节点的颜色. 如果 $v$ 被合并到 $u$, 最后有 $color(u)=R_2$, 那么自然得到 $color(v)=R_2$.
+
+Coalescing 还会改变 degree. 当 $v$ 的边转移给 $u$ 后, 某些邻居可能新增与 $u$ 的 interference, 某些边则已经存在, 不需要重复计算. 合并后的 $u$ 也可能从 low-degree 变成 high-degree, 因而从一个适合 freeze 的节点转移到 spill candidate 集合. 这说明 coalescing 和 simplify 不能各自独立运行一次, 两者需要反复交替.
+
+## 4 为什么需要 Freeze
+
+假设一个节点 $v$ 满足 $degree(v)<K$. 按照第四部分的 simplify 原理, 它可以安全移出图. 但如果 $v$ 还参与 COPY, 立刻 simplify 会错过后面进行 coalescing 的机会.
+
+因此 IRC 会进一步区分 low-degree nodes. 没有活跃 move relation 的 low-degree node 可以进入 `simplifyWorklist`; 仍然 move-related 的 low-degree node进入 `freezeWorklist`. 后者暂时保留在图中, 让算法先尝试处理它关联的 COPY.
+
+这里的 move-related 可以理解成节点仍然参与某个尚未决定命运的 move. 假设 $moveList(v)$ 保存所有与 $v$ 有关的 moves, 那么只要其中还有 move 位于 `worklistMoves` 或 `activeMoves`, $v$ 就仍然被视为 move-related.
+
+问题是, 某些 COPY 经过 conservative test 后暂时不能 coalesce, 图里又可能没有别的可合并 move. 如果所有 low-degree nodes 都因为 move relation 停在 `freezeWorklist`, allocator 就无法继续 simplify.
+
+Freeze 的作用是在这种时候主动放弃某些 COPY 的 coalescing 机会.
+
+假设有:
+
+```text
+v = COPY u
+```
+
+当前无法安全 coalesce, 但 $v$ 是 low-degree node. Freeze $v$ 时, allocator 会把与 $v$ 相关的尚未处理 move 标记为 frozen. 这相当于接受 "这条 COPY 以后可能保留下来" 这个结果. 一旦这些 move 不再被视为待 coalesce, $v$ 就可能变成 non-move-related node, 随后进入 `simplifyWorklist`.
+
+Freeze 不会给两个节点建立 interference edge, 也没有强制它们最终使用不同颜色. 它只是停止主动追求这条 move 的合并. 如果最后 coloring 恰好让两端获得同一个寄存器, COPY 仍然有可能被后续清理掉. IRC 在这里放弃的是 coalescing 的算法保证和继续等待的成本.
+
+这样一来, low-degree node 的状态大致会在两种工作集合之间转换:
+
+```text
+low degree + not move-related
+        |
+        v
+simplifyWorklist
+
+low degree + move-related
+        |
+        v
+freezeWorklist
+```
+
+而 $degree\ge K$ 的普通节点通常进入 `spillWorklist`. 名字里的 `spillWorklist` 同样不意味着这些节点已经确定要 spill, 它们仍然只是 potential spill candidates.
+
+## 5 Iterated Register Coalescing 的整体流程
+
+IRC 的核心设计是让 Simplify, Coalesce, Freeze 和 SelectSpill 在同一个循环中不断改变图的状态. 一次 coalescing 可能降低或提高某些节点的处理优先级, simplify 会降低邻居 degree, degree 下降以后原来不安全的 coalescing 又可能变得安全, freeze 又可以解除 move relation 对 simplify 的阻塞. 所以算法不断在这些操作之间循环, 直到所有节点都被移出当前图.
+
+一个典型的顶层结构可以写成:
+
+```text
+Build
+MakeWorkList
+
+while simplifyWorklist is not empty
+   or worklistMoves is not empty
+   or freezeWorklist is not empty
+   or spillWorklist is not empty:
+
+    if simplifyWorklist is not empty:
+        Simplify()
+
+    else if worklistMoves is not empty:
+        Coalesce()
+
+    else if freezeWorklist is not empty:
+        Freeze()
+
+    else:
+        SelectSpill()
+
+AssignColors()
+
+if spilledNodes is not empty:
+    RewriteProgram()
+    restart allocation
+```
+
+这个优先顺序体现了算法的偏好. 有安全可 simplify 的节点时先 simplify; 有待处理 COPY 时尝试 coalesce; coalescing 暂时无法推进时才 freeze; 只有这些方法都不能推进图时才选择 potential spill.
+
+`Simplify()` 和第四部分基本一致. 从 `simplifyWorklist` 中拿一个节点 $n$, 压入 `selectStack`, 然后删除它对当前图的影响. 每个邻居的 degree 减 1. 如果某个邻居原本恰好满足 $degree=K$, 减少以后变成 $K-1$, 它从 high-degree 变成 low-degree, 这可能改变它所在的 worklist.
+
+如果这个邻居已经没有待处理 moves, 可以进入 `simplifyWorklist`. 如果仍然 move-related, 则进入 `freezeWorklist`. 同时 degree 的下降还可能让与它关联的 `activeMoves` 重新满足 conservative coalescing 条件, 因此这些 moves 可以重新进入待尝试状态.
+
+`SelectSpill()` 则从 `spillWorklist` 中根据 spill heuristic 选择一个节点 $m$. 它不会立刻生成 load/store, 只是把 $m$ 移到 simplify 流程中, 相当于接受它可能在 AssignColors 阶段失败. 这就是上一部分讲过的 optimistic coloring 在 IRC 中的体现.
+
+## 6 一个 move 在 IRC 中会经历哪些状态
+
+IRC 不只维护节点 worklists, move 本身也有状态.
+
+新构建出来, 还等待尝试 coalescing 的 move 放在 `worklistMoves`. 当 `Coalesce()` 取出一条 move $x\leftarrow y$ 时, allocator 先通过 alias 找到当前真正的两个代表节点 $u=GetAlias(x)$ 和 $v=GetAlias(y)$.
+
+如果 $u=v$, 说明之前的其他 coalescing 已经间接把两端合成了同一个节点. 这条 move 可以直接放入 `coalescedMoves`.
+
+如果 $u$ 和 $v$ 已经存在 interference edge, 那么它们不能共享颜色, move 进入 `constrainedMoves`. 这种 COPY 无法通过 coalescing 消失.
+
+如果两端没有 interference, 并且通过 George 或 Briggs 的 conservative test, allocator 执行 `Combine(u,v)`, move 同样进入 `coalescedMoves`.
+
+还有一种情况是两端目前没有 interference, 但 conservative test 暂时失败. 这并不意味着以后永远不能合并. 随着 simplify 删除邻居, degree 可能下降, 原来危险的 coalescing 可能变得安全. 因此 IRC 把这种 move 放进 `activeMoves`, 暂时等待图结构变化.
+
+最后, Freeze 主动放弃某条 move 的 coalescing 尝试时, 它会进入 `frozenMoves`.
+
+于是这些集合表达了 move 的不同命运:
+
+|集合|含义|
+|---|---|
+|`worklistMoves`|等待尝试 coalescing|
+|`activeMoves`|当前不安全, 等待图变化|
+|`coalescedMoves`|两端已经合并|
+|`constrainedMoves`|两端存在硬冲突, 无法合并|
+|`frozenMoves`|allocator 主动停止尝试合并|
+
+这里 `activeMoves` 的存在很有意义. 如果一次 conservative test 失败就永久放弃 move, simplify 后降低 degree 所带来的新机会就利用不上. IRC 中的 "Iterated" 很大程度上就体现在这种重新激活上.
+
+## 7 节点 worklist 如何随着 degree 和 move 状态变化
+
+节点也不是从初始化开始就固定属于某个集合. Build 完成后, allocator 通常根据 degree 和 move relation 把普通节点分成三类.
+
+如果 $degree(n)\ge K$, 节点进入 `spillWorklist`. 如果 $degree(n)<K$ 且仍然 move-related, 进入 `freezeWorklist`. 如果 $degree(n)<K$ 且已经没有待处理 moves, 进入 `simplifyWorklist`.
+
+假设一个节点 $n$ 最初满足 $degree(n)=K$, 所以位于 `spillWorklist`. 某个邻居被 simplify 后, $degree(n)$ 下降到 $K-1$. 这时 $n$ 已经具有低 degree. 如果它还参与待处理 COPY, 它会转入 `freezeWorklist`; 如果没有, 就转入 `simplifyWorklist`.
+
+反过来, coalescing 也可能让一个 low-degree node 的邻居集合扩大. 如果合并后的代表节点 degree 升到 $K$ 或以上, 它可能重新进入 high-degree 状态.
+
+因此 IRC 的数据结构本质上维护着一个不断变化的分类系统. Degree 描述 coloring 风险, move-related 描述 coalescing 机会, 两者共同决定一个节点当前应该被 simplify, freeze, 还是暂时作为 spill candidate 保留.
+
+## 8 用一个小例子走一遍 Simplify, Coalesce 和 Freeze
+
+假设 $K=3$, 当前有节点 $a,b,c,d,e$, 并存在两条 COPY:
+
+```text
+b = COPY a
+e = COPY d
+```
+
+冲突关系为:
+
+```text
+a -- c
+b -- c
+b -- d
+c -- d
+d -- e
+```
+
+同时 $a-b$ 没有 interference, 因此第一条 COPY 有机会 coalesce. $d-e$ 已经存在 interference edge, 所以第二条 COPY 无法让两端共享颜色.
+
+开始时, `b = COPY a` 位于 `worklistMoves`. allocator 检查 $a$ 和 $b$ 的邻居集合. 假设合并通过 conservative test, 就把 $a$ 和 $b$ 合并成节点 $ab$.
+
+原来的关系:
+
+```text
+a -- c
+b -- c
+b -- d
+```
+
+合并后变成:
+
+```text
+ab -- c
+ab -- d
+```
+
+而 `b = COPY a` 进入 `coalescedMoves`.
+
+对于 `e = COPY d`, 由于 $d-e$ 已经有 interference edge, 这条 move 进入 `constrainedMoves`. allocator 不会再尝试合并它们. 如果 $e$ 此时 degree 很低并且没有其他活跃 move, 它可以进入 `simplifyWorklist`.
+
+随着 $e$ 被 simplify, $d$ 的 degree 下降. 如果 $d$ 原本处于 high-degree 状态, 它可能因此跌到 $K$ 以下, 从 `spillWorklist` 转移到 `simplifyWorklist` 或 `freezeWorklist`.
+
+再考虑另一种情况. 假设 $a-b$ 的 coalescing 第一次因为 high-degree neighbors 太多而没有通过 conservative test, move 会进入 `activeMoves`. 随后若若干邻居被 simplify, $a$ 或 $b$ 周围的 degree 降低, allocator 会重新启用这条 move. 第二次检查时它可能已经满足 Briggs criterion, 于是完成 coalescing.
+
+如果图已经无法继续 simplify, active moves 也没有变得安全, 但存在低 degree 且 move-related 的节点, Freeze 会选择其中一个节点, 把相关 moves 转入 `frozenMoves`. 该节点失去 move-related 状态以后就能进入 simplify. 这样算法不会因为追求 COPY elimination 而卡在原地.
+
+## 9 AssignColors 与 coalesced nodes
+
+IRC 的主循环结束后, 普通节点都已经被压入 `selectStack`, 或者通过 coalescing 成为了另一个节点的 alias. 接下来的 `AssignColors()` 与第四部分的 Select 类似: 依次从栈顶弹出节点, 查看已经着色的邻居占用了哪些物理寄存器, 从合法颜色中选择一个剩余颜色.
+
+对于普通节点 $n$, 可以先构造被邻居阻塞的颜色集合 $Forbidden(n)$, 然后计算 $Available(n)=LegalColors(n)-Forbidden(n)$. 如果 $Available(n)$ 为空, $n$ 加入 `spilledNodes`; 否则从其中选择一个物理寄存器. 如果存在 register preference 或 move hint, allocator 可以优先选择有助于删除 COPY 的颜色.
+
+所有 ordinary nodes 处理完以后, coalesced nodes 通过 alias 继承颜色. 如果 $alias(v)=u$, 那么令 $color(v)=color(u)$. 所以一条成功 coalesced 的 `v = COPY u` 最终天然满足 $color(v)=color(u)$.
+
+如果 `spilledNodes` 非空, allocator 会进行 spill rewrite, 重新计算 liveness 和 interference, 再重新运行整个算法. Spill rewrite 产生的新 temporaries 和新的 COPY 也会重新进入下一轮 IRC. 因此 coalescing 和 spilling 之间也存在反馈: 某次 spill 可能拆短 live range, 让下一轮的 COPY 更容易合并; coalescing 也可能增加 pressure, 最终使某个节点发生 spill.
+
+## 10 IRC 中的数据结构
+
+把算法实现出来时, 需要同时保存图结构, 节点状态和 move 状态. `adjList[n]` 或 `adjSet` 保存 interference relation, `degree[n]` 保存当前简化图中的 degree, `moveList[n]` 保存与节点相关的 COPY, `alias[n]` 记录 coalescing 后的代表节点, `color[n]` 保存最终物理寄存器.
+
+节点状态通常包括 `simplifyWorklist`, `freezeWorklist`, `spillWorklist`, `selectStack`, `coalescedNodes`, `coloredNodes` 和 `spilledNodes`. 前三个表示当前仍在图中的节点分别等待哪种操作; `selectStack` 保存已经 simplify 或 optimistic remove 的节点; `coalescedNodes` 保存被其他节点吸收的节点; 最后的 `coloredNodes` 和 `spilledNodes` 是 AssignColors 后的结果.
+
+这套数据结构看起来繁琐, 原因在于 IRC 同时维护两个会互相影响的过程. Coloring 希望不断删除低 degree nodes, coalescing 希望暂时保留 move-related nodes 等待合并机会. Degree 的变化会改变 coalescing 条件, move 的冻结或合并又会改变节点是否能够 simplify. Worklist 的作用就是明确记录每个对象当前处于哪一个阶段, 避免每次都重新扫描整张图.
+
+从算法结构上看, IRC 已经比单纯的 Chaitin coloring 完整很多. Chaitin 的主要循环围绕 simplify 和 spill 展开, IRC 把 COPY optimization 纳入同一个状态机, 形成:
+
+```text
+               +------------------+
+               |                  |
+               v                  |
+           Simplify               |
+               |                  |
+               v                  |
+            Coalesce ------------>+
+               |
+               v
+             Freeze
+               |
+               v
+          SelectSpill
+               |
+               +------> Simplify
+
+main loop ends
+       |
+       v
+ AssignColors
+       |
+       v
+ Rewrite if spilled
+```
+
+实际执行并不会严格按图中的固定环路逐项走一遍. 每次循环都会根据当前 worklist 是否为空选择能够推进的操作, 图和 move 状态随之变化.
+
+## 11 Coalescing 的收益和代价
+
+Coalescing 最直接的收益是减少 register-to-register move. 对 CPU 来说, 某些 move 可能在 rename 阶段成本很低, 某些甚至可以被硬件消除, 但它们仍然可能占用 decode/issue bandwidth, 增加 code size, 并影响调度. 对一些目标机器或特殊 register classes 来说, move 的成本还会更高.
+
+另一方面, coalescing 会延长或合并 live ranges. 原本两个不同时期占用寄存器的值一旦被视为一个更大的 allocation object, 其邻居集合可能增大, register pressure 也可能在某些区域更难处理. 所以 "消除更多 COPY" 和 "避免更多 spill" 并不总是同一个方向.
+
+这也是 conservative coalescing 的基本取舍. Aggressive coalescing 更愿意删除 COPY, 可能付出更多 spill 风险; conservative coalescing 会保留一些 move, 换取更稳定的 colorability. IRC 进一步让这个选择随着 simplify 动态变化, 第一次不安全的 move 可以进入 `activeMoves`, 等 degree 降低后重新尝试.
+
+现代工业 allocator 未必直接实现完整的 textbook IRC 状态机, 但这里形成的几个观念会一直保留下去: COPY 可以形成 register preference, coalescing 会改变 live range 和 interference, 合并需要考虑 pressure, 某些 move 可以推迟处理, 也可以在必要时放弃. 后面看 LLVM RegisterCoalescer, LLVM Greedy 的 register hints, 以及 GCC IRA 的 copy cost 时, 都能看到这些思想的延续.
+
+下一部分进入 Linear Scan Register Allocation. 到那里我们会暂时放下显式 interference graph, 改用 live interval 的线性顺序处理寄存器竞争. 这样可以直接比较两种经典视角: graph coloring 主要保存 "谁和谁冲突", linear scan 则主要保存 "每个值在什么时候存活".
diff --git a/content/notes/register-allocation/cpu-constraints/index.md b/content/notes/register-allocation/cpu-constraints/index.md
new file mode 100644
index 0000000..7f99f15
--- /dev/null
+++ b/content/notes/register-allocation/cpu-constraints/index.md
@@ -0,0 +1,243 @@
+---
+title: "第八部分: 真实 CPU 体系结构约束"
+slug: "cpu-constraints"
+lang: zh-Hans
+series: register-allocation
+weight: 90
+created: 2026-08-18
+updated: 2026-09-09
+license: CC-BY-SA-4.0
+---
+
+前面的 graph coloring 和 Linear Scan 都从一个相对理想化的模型出发. 假设机器有 $K$ 个物理寄存器, 每个 virtual register 都可以从这 $K$ 个寄存器中任选一个, 相邻 live ranges 只需要保证颜色不同. 这个模型足以解释 interference, spilling, coalescing 和 splitting 的基本原理, 但真实 CPU 的 register file 很少具有这种完全对称的结构.
+
+一条机器指令可能只接受某个 register class, 某些 operands 必须使用固定 physical register, 一组不同名字的寄存器可能实际覆盖同一片硬件存储, function call 会破坏一部分 registers, two-address instruction 又可能要求 input 和 output 使用同一个 register. 此外, 某些 operand 还需要连续寄存器, register pair, 特定对齐或特殊 subregister.
+
+因此工业 register allocation 面对的已经不只是 "$K$ 个等价颜色上的 graph coloring". 每个 allocation object 都可能拥有自己的合法寄存器集合, 不同 physical registers 之间存在 overlap, 某些程序位置还会临时禁止一部分 registers. Register allocator 必须把这些 ISA 和 ABI 约束一起纳入 allocation.
+
+## 1 Register Class 与不对称的颜色集合
+
+最直接的机器约束来自 register class.
+
+以典型 CPU 为例, integer arithmetic 使用 general-purpose registers, floating-point 和 SIMD instructions 使用 vector registers. 一个整数 virtual register 不能因为 GPR 不够用就随意放到 vector register 中, 反过来也一样. 即使底层硬件拥有很多 register bits, ISA 允许某条 instruction 使用哪些寄存器仍然受到 operand encoding 和 instruction semantics 限制.
+
+因此对 virtual register $v$, allocator 实际面对的不是统一的 $K$ 个颜色, 而是一个合法 physical register 集合 $C(v)$. Assignment 至少必须满足 $color(v)\in C(v)$.
+
+例如某台抽象机器有:
+
+```text
+GPR:
+R0 R1 R2 R3
+
+Vector:
+V0 V1 V2 V3 V4 V5 V6 V7
+```
+
+如果 `v1` 是 integer value, 它的 candidate set 可能是 ${R0,R1,R2,R3}$. 如果 `v2` 是 vector value, candidate set 则是 ${V0,\ldots,V7}$. 即使 `v1` 和 `v2` 在同一个程序位置同时 live, 它们也未必竞争同一份物理资源, 所以简单计算 "总共有多少 live values" 已经不能准确表示 register pressure.
+
+更复杂的情况是 register classes 可以部分重叠. 某些 instruction 接受一个较大的 class, 另一些 instruction 只接受其中的子集. 假设机器有八个 general-purpose registers, 但某条特殊 instruction 只能编码 `R0-R3`, 那么参与这条 instruction 的 operand 在那个位置附近实际受到更严格的约束.
+
+这会改变 graph coloring 的数学模型. 经典 $K$-coloring 假设所有节点共享同一个颜色集合. 每个节点拥有自己的 candidate set 后, 问题更接近 list coloring. 两个 degree 完全相同的 virtual registers, 如果一个有 12 个候选物理寄存器, 另一个只有 2 个, allocation difficulty 显然不同.
+
+所以工业 allocator 在判断 "这个 live range 难不难分配" 时, degree 或 live range overlap 只是其中一部分. Register class 大小, 当前被占用的 candidate registers, target instruction constraints 和 register hints 都会影响实际可分配性.
+
+Register class 还会改变 spill 的意义. 某个大 register class 仍有很多空闲资源, 并不能帮助另一个已经超压的小 register class. Register pressure 因而通常需要按照不同 register resources 分别追踪.
+
+## 2 Fixed Register, Pre-colored Value 与 Calling Convention
+
+有些机器值在进入 allocator 之前就已经确定了 physical register. Graph-coloring 文献通常把这种节点称为 pre-colored node.
+
+Calling convention 是最常见的来源之一. 函数参数, 返回值和某些特殊 runtime state 会按照 ABI 放在指定 registers 中. 例如一个返回值可能要求位于某个固定 GPR, allocator 不能随意把最终 return operand 放到另外一个寄存器然后假设调用者能够理解.
+
+机器指令自身也会产生 fixed-register constraints. x86 的一些指令是典型例子. 整数除法会隐式使用固定的 accumulator 和 high-half registers. 传统 variable shift 指令的 shift count 可能要求位于 `CL`. 这些 instructions 让某个程序位置上的 register availability 突然发生变化.
+
+假设有一个 live range `v` 跨过某条必须使用 `R0` 的 instruction:
+
+```text
+v:
+|------------------------------|
+
+                  fixed R0 use
+                       |
+                       v
+-----------------------X--------
+```
+
+如果 `v` 自己被分配到 `R0`, allocator 就必须确认它和这个 fixed use 是否能够合法共存. 如果固定指令会破坏 `R0`, `v` 就不能无条件跨过该位置继续保存在 `R0`. 一种方案是选择其他 physical register, 另一种方案是在固定使用周围 split `v`.
+
+Linear Scan 可以把这种约束表示成 physical register 的 fixed interval. Graph-based allocator 可以把 physical register 视为 pre-colored object, 再建立相应 interference. 工业实现的数据结构不同, 表达的机器事实是一样的: 某些物理资源在特定位置已经被机器语义占用.
+
+Function call 把这种问题扩展到了 ABI 层面. 大多数 ABI 会把 registers 分成 caller-saved 和 callee-saved 两类. Caller-saved register 可以被被调用函数自由破坏, 因此一个跨 call live 的 value 如果放在这种 register 中, 当前函数需要在 call 周围保存它, 或通过其他方式保证值不丢失.
+
+```text
+v = ...
+...
+call foo
+...
+use v
+```
+
+如果 `v` 被分配到一个会被 `foo` 破坏的 caller-saved register, allocator 可能需要:
+
+```text
+store [slot], R0
+call foo
+load R0, [slot]
+```
+
+如果把 `v` 放进 callee-saved register, 它可以自然跨越 call, 但当前函数一旦使用这个 callee-saved register, 通常就需要在 prologue 和 epilogue 中保存和恢复它.
+
+因此 caller-saved 和 callee-saved 并不是简单的 "坏寄存器" 和 "好寄存器". 一个很短且不跨 call 的 value 使用 caller-saved register 通常很自然. 一个跨越多个 calls 的长 live range 可能更适合 callee-saved register. 如果整个函数只为了一个很冷的 value 使用一次新的 callee-saved register, prologue/epilogue 成本又可能超过局部 spill 的成本.
+
+Physical registers 因而开始具有不同的 assignment cost. 从这个阶段开始, register allocation 已经不能只问 "哪个颜色合法", 还要问 "合法颜色中哪个更便宜".
+
+## 3 Subregister 与 Physical Register Aliasing
+
+真实 register file 中另一个会破坏简单 coloring 模型的结构是 aliasing.
+
+x86 的 `RAX`, `EAX`, `AX` 和 `AL` 是最熟悉的例子. 它们拥有不同的 architectural names 和 operand widths, 但覆盖的是同一个物理寄存器的不同部分:
+
+```text
+RAX: |------------------------------- 64 bits -------------------------------|
+
+EAX:                                 |------------- 32 bits -----------------|
+
+ AX:                                                 |------ 16 bits --------|
+
+ AL:                                                        |-- 8 bits -----|
+```
+
+所以如果一个 live value 占用了 `RAX`, 另一个同时 live 的 value就不能独立占用 `EAX` 或 `AX`. 从寄存器分配角度看, 这些 names 不是互不相关的颜色.
+
+AArch64 的 `X0` 和 `W0` 也存在类似的资源关系. SIMD ISA 中还可能出现不同宽度 register views 共享相同底层存储的情况. 对 x86 AVX 系列来说, `XMM0`, `YMM0` 和 `ZMM0` 也对应不同宽度的重叠 register views.
+
+因此物理冲突更准确的条件不是 "两个 values 是否被分配到同一个 register name", 而是它们所占用的底层 register resources 是否 overlap. 如果用 $Units(r)$ 表示 physical register $r$ 占用的硬件资源集合, 两个同时 live 的 values $u$ 和 $v$ 至少需要满足 $Units(color(u))\cap Units(color(v))=\varnothing$.
+
+这会让 register allocation 的 interference 检查明显复杂起来. 假设一个 value 可以分配到 `EAX`, 另一个 value 可以分配到 `RAX`. 两个 candidate names 不同, 但不能同时使用. allocator 因而通常需要维护 physical-register overlap relation 或更细的 register units, 而不是只比较 register number.
+
+Subregister 还带来读写语义问题. 以 x86 为例, 写 `EAX` 会把 `RAX` 的高 32 位清零, 因而这种写操作不仅覆盖低 32 位. 写更小的 subregister 又可能保留其余 bits. 所以 "哪些 names alias" 和 "一次 definition 实际覆盖哪些 bits" 并不总是同一个简单问题.
+
+后端通常通过 target-specific register description 表达这些关系. Register allocator需要知道哪些 assignments 会发生物理冲突, 后续 liveness 和 dead-def analysis 还要理解不同 subregister definitions 对 super-register value 的影响.
+
+这一层机器语义也说明了为什么 source-level type width 不能直接决定一个 value 占用多少 register resource. 一个 32-bit value放在 `EAX` 时, 对 register allocation 来说通常已经占据了 `RAX` 对应的那份 GPR 资源, 不存在另一个同时 live 的 32-bit value 可以偷偷使用 `RAX` 的高半部分.
+
+## 4 Tied Operand, Two-address Instruction 与 Early-clobber
+
+前面讨论的 interference 大多来自不同 values 的生命周期重叠. ISA 还可以直接规定几个 operands 之间必须满足某种位置关系.
+
+Two-address instruction 是典型情况. 一个三地址抽象操作可以写成:
+
+```text
+v3 = ADD v1, v2
+```
+
+某些 ISA 或具体 instruction encoding 则要求 destination 与一个 source 使用同一个 physical register:
+
+```text
+R0 = ADD R0, R1
+```
+
+在 machine IR 中, 这通常表现为 output operand 和某个 input operand tied. 如果进入这条 instruction 时有 `v1` 和 `v2`, 结果是 `v3`, allocator希望满足类似 $color(v_3)=color(v_1)$ 的 tied constraint.
+
+如果 `v1` 在 instruction 后仍然 live, 问题就会出现. Destination 覆盖 `v1` 以后, 后续 use 会失去旧值. Compiler 可能需要提前复制:
+
+```text
+tmp = COPY v1
+tmp = ADD tmp, v2
+...
+use v1
+```
+
+这样 `tmp` 可以和 result 共用 physical register, 原来的 `v1` 继续保存在别处.
+
+所以 tied operand 与第五部分的普通 coalescing preference 不完全一样. 普通 COPY 两端希望同色, 失败后通常只是留下 move. Tied operand 是机器 instruction legality 的一部分, allocator 或之前的 lowering 必须保证最终满足这个关系.
+
+Early-clobber 则产生相反方向的额外 interference. 对普通 instruction, allocator 往往可以认为所有 inputs 先被读取, output 随后才写入. 如果一个 input 在这条 instruction 后死亡, destination 有时可以复用这个 input 的 register.
+
+Early-clobber operand 会更早写入 destination, 早到某些 inputs 尚未全部读取. 于是 destination 不能和那些 inputs 使用同一个 physical register, 即使从普通 instruction-boundary liveness 看, input 的生命周期似乎正好在这里结束.
+
+可以用一个抽象 instruction 表示:
+
+```text
+early-def dst, use src1, use src2
+```
+
+假设 `dst` 在 instruction 执行早期就被覆盖, 而 `src2` 到更晚的阶段才读取. 那么 `dst` 和 `src2` 必须分配不同的物理资源. 这种冲突来自 instruction 内部的 operand timing, 单纯使用 "指令前 live / 指令后 live" 两个边界不足以表达.
+
+这也是前面提到 machine position 需要细分的原因. Register allocator必须能够理解 ordinary use, early use, normal def, early-clobber def 等不同时间点, 否则会产生在抽象 liveness 上看似合法, 实际 ISA 无法执行的 assignment.
+
+Implicit use 和 implicit def 也属于同一类机器约束. 某些 instruction encoding 没有把 register 明写在 operand list 中, 但架构语义仍然会读取或修改它. Register allocator必须把这些隐含 effects 纳入 physical-register liveness.
+
+## 5 Register Pair, Multiple-register Operand 与组合资源
+
+简单 coloring 还假设一个 value 对应一个 physical register. 某些机器操作需要同时占据多个寄存器.
+
+一个常见来源是 value width 超过单个 register width. 在 32-bit target 上处理 64-bit 或更宽整数时, 一个 logical value 可能被拆成多个 machine values. 某些 ISA 或 calling convention 还会要求它们使用特定 pair 或具有特定排列关系.
+
+另一些 instruction 本身就操作 register pairs. x86 的某些 multiply/divide 语义会把结果或输入分布在固定的一对 registers 中. 其他架构还可能要求一组 operands 使用连续 registers, even-odd pair 或满足一定 alignment.
+
+这时 allocator 的 assignment unit 不再总是单个颜色. 假设某个 virtual object 需要一对 consecutive registers, 候选方案可能是:
+
+```text
+(R0, R1)
+(R2, R3)
+(R4, R5)
+```
+
+即使当前还有两个空闲寄存器, 如果它们是 `R1` 和 `R4`, 也无法满足这个 allocation object 的要求.
+
+组合资源使 register pressure 更加具有结构性. "还有几个寄存器空闲" 不足以判断一个 tuple 是否能放下, allocator还需要知道空闲资源的具体形状. 这也是某些 target 上 fragmentation 会影响 allocation 的原因: 总空闲容量足够, 合法组合却已经不存在.
+
+CPU 上这种问题的复杂程度因 ISA 而异. 到 GPU 部分还会看到更明显的 tuple constraints 和不同 register files, 但基本思想相同: allocator分配的不一定是单个独立 register name, 也可能是一组必须同时满足结构要求的硬件资源.
+
+## 6 Machine Constraint 如何改变 Spill 与 Split
+
+当 allocation 失败时, spilling 也必须服从 target ISA. 把一个 value 标记为 "memory" 并不能自动解决所有问题, 因为很多 machine instructions 不允许任意 operand直接来自 memory.
+
+假设:
+
+```text
+v3 = ADD v1, v2
+```
+
+如果目标 instruction 要求两个 inputs 都在 registers 中, spill `v1` 后仍然必须在 use 前产生 reload:
+
+```text
+t1 = load [slot]
+v3 = ADD t1, v2
+```
+
+`t1` 又形成一个新的短 live range, 仍然需要 physical register. 所以 spill 没有消除寄存器需求, 它只是把一个长时间占据资源的 value 转换成 use 附近的短期需求.
+
+某些 ISA 允许 memory operand. 例如一条 arithmetic instruction 可能允许一个 source 直接来自内存. 这时 spill rewrite 有机会把 reload folding 进原来的 machine instruction, 减少显式 load. 这种收益完全依赖 instruction encoding 和 target cost model.
+
+Live range splitting 同样可以围绕机器约束进行. 如果一个长 live range 大部分时间可以使用任意 GPR, 但在一个局部区域必须进入特定 register subset, allocator 可以在这个区域附近 split:
+
+```text
+general class        restricted class        general class
+|----------------|----------------------|------------------|
+```
+
+不同 fragments 分别选择最合适的 physical register, 边界之间通过 COPY 连接. 如果 allocation 恰好让相邻 fragments 得到同一个 register, COPY 可以被消除.
+
+Function call 周围也是常见 split point. 一个 value 在 call 之前和之后都需要使用, 但跨 call 保存在某个 register 中代价很高, allocator可以把 live range 在 call 附近切开. 前一段使用 caller-saved register, call 时临时进入 stack, 后一段 reload 到另一个 register. 对只跨少数 calls 的 value, 这种局部决策可能比占用 callee-saved register 整个函数更便宜.
+
+Rematerialization 也会受到 target instruction cost 影响. 如果某个 value 是一个便宜的 immediate 或简单地址计算, 在 use 处重新生成可能比真正 reload 更合适. 对 register allocator来说, 这相当于把 memory spill cost 替换成一段 machine computation cost.
+
+所以 machine constraints 并没有只增加 "哪些寄存器可以用" 这一层限制. 它们还会改变 spill placement, split point, copy elimination 和 rematerialization 的成本模型.
+
+## 7 从 K-coloring 到真实 Machine Allocation
+
+经过这些约束以后, textbook coloring 模型仍然有价值, 但需要重新理解其中的 "颜色".
+
+对简单 graph coloring, 每个节点从同一集合 ${R_0,\ldots,R_{K-1}}$ 中选择一个颜色, interference edge $(u,v)$ 要求 $color(u)\neq color(v)$.
+
+真实机器上, 每个 value $v$ 有自己的 candidate set $C(v)$; physical registers 之间可能 alias; 某些 instruction positions 存在 fixed uses 和 clobbers; tied operands 会要求两个 values 获得相同或相关 assignment; early-clobber 又会制造额外的局部 interference; 某些 objects 甚至需要一次分配多个 physical registers.
+
+因此一条 interference edge 最终表达的机器条件更接近: 两个同时 live 的 allocation objects 不能占用重叠的 physical register resources. 如果两个 register names 不同但底层 units overlap, assignment 仍然非法.
+
+Graph coloring allocator 可以通过 register classes, pre-colored nodes, special interference 和 target-specific constraints 扩展这个模型. Linear Scan 可以通过 candidate sets, fixed intervals 和位置约束实现类似功能. Priority-based allocator还可以在遇到冲突时尝试 eviction, splitting 和重新排队.
+
+这也是接下来进入 LLVM 和 GCC 工业实现时需要注意的. LLVM Greedy 不会把整个问题真的展开成一张 textbook interference graph, 但它仍然必须回答本章的所有机器问题: virtual register 属于哪个 register class, 某个 physical register 是否与当前 live interval overlap, subregister alias 如何计算, fixed interference 在哪里, tied operand 如何满足, call clobber 如何建模, splitting 后 fragment 可以使用哪些 registers.
+
+GCC IRA 和 LRA 则把其中一部分问题分成两个阶段. IRA 更侧重全局 allocation 和 register pressure, LRA 更靠近机器 instruction constraints, reload 和最终 hard-register assignment. 这种架构上的差异会在下一部分具体展开.
diff --git a/content/notes/register-allocation/fundamentals/index.md b/content/notes/register-allocation/fundamentals/index.md
new file mode 100644
index 0000000..accb6e6
--- /dev/null
+++ b/content/notes/register-allocation/fundamentals/index.md
@@ -0,0 +1,825 @@
+---
+title: "第一部分：寄存器分配问题与基本模型"
+slug: "fundamentals"
+lang: zh-Hans
+series: register-allocation
+weight: 20
+created: 2026-08-09
+updated: 2026-09-09
+license: CC-BY-SA-4.0
+---
+
+## 1. 从虚拟寄存器到物理寄存器
+
+编译器进入后端以后,  会逐渐把源语言里的表达式, 变量和控制流降低成越来越接近目标机器的形式. 到了指令选择之后, 程序通常已经表现为机器指令或者准机器指令, 但这些指令里使用的寄存器仍然常常是**虚拟寄存器**.
+
+例如, 一个中间结果可能被表示成:
+
+```text
+%v1 = ADD %v2,   %v3
+%v4 = MUL %v1,   %v5
+%v6 = SUB %v4,   %v1
+```
+
+这里的 `%v1`, `%v2`, `%v3` 等都可以理解为虚拟寄存器. 编译器可以创建任意多个这样的名字, 因此在这一层上, 程序暂时不需要关心目标 CPU 到底只有多少个真实寄存器.
+
+这种设计很重要. 假如编译器从很早的阶段开始就强迫所有中间值使用真实寄存器, 那么几乎每一次代码移动, 公共子表达式消除, 循环优化, 指令选择都会受到具体寄存器数量的牵制. 虚拟寄存器提供了一层抽象, 使前面的优化阶段可以先专注于数据依赖和程序语义.
+
+到了真正生成机器代码的时候, 这层抽象必须被消除.
+
+假设目标机器只有:
+
+```text
+R0
+R1
+R2
+R3
+```
+
+四个可以使用的通用寄存器, 那么所有仍然存活的虚拟寄存器最终都必须找到实际存储位置. 某些虚拟寄存器会被映射到 `R0`, 某些会映射到 `R1`. 两个生命周期互不重叠的虚拟寄存器还可以在不同时间复用同一个物理寄存器.
+
+因此,  寄存器分配最核心的问题总结起来就是
+
+> 给程序中的值安排物理寄存器,  使所有机器约束得到满足, 并尽可能降低由于寄存器不足带来的执行代价.
+
+---
+
+## 2. Register Allocation, Register Assignment 与 Spilling
+
+"Register Allocation" 这个词在不同文献和编译器中使用得并不完全一致.
+
+从理论上细分时, 经常会区分两个问题.
+
+**Register Allocation** 决定某个值在某段生命周期中是否应该驻留在寄存器里.
+
+**Register Assignment** 决定已经确定要放进寄存器的值具体使用哪一个物理寄存器.
+
+例如:
+
+```text
+v1 → R2
+v2 → R0
+v3 → R2
+```
+
+这里 `v1` 和 `v3` 可以同时使用 `R2`, 前提是它们的生命周期没有发生重叠.
+
+工业编译器中,  "register allocation" 往往用作整个过程的总称, 其中通常会同时包含分配, 具体寄存器选择, spill, live range splitting, coalescing, rematerialization 等步骤. 因此以后我们讨论 LLVM 的 Greedy Register Allocator 或 GCC IRA 时, 使用的基本也是这个宽泛含义.
+
+真正困难的情况出现在同时需要保留的值数量超过可用物理寄存器数量时.
+
+假设某段程序中同时存活"
+
+```text
+a
+b
+c
+d
+e
+```
+
+而当前只能使用:
+
+```text
+R0
+R1
+R2
+R3
+```
+
+那么五个值不可能全部同时驻留在这四个寄存器中.
+
+编译器需要让其中至少一个值暂时离开寄存器. 最经典的方法是把它保存到栈上的一个位置中:
+
+```text
+store [spill_slot],   R1
+```
+
+稍后再次需要这个值时, 再把它读取回来:
+
+```text
+load R1,   [spill_slot]
+```
+
+这个过程称为 **spilling**.
+
+被写入内存的值通常称为 spilled value, 相应的内存位置常称为 spill slot.
+
+因此, 寄存器分配从一开始就包含两个彼此相关的决定:
+
+```text
+哪些值值得留在寄存器里?
+
+留下来的值应该使用哪个寄存器?
+```
+
+后面我们会看到, 很多经典算法虽然在形式上强调 "给冲突图着色", 实际的算法质量往往很大程度取决于第二层以外的问题: 什么时候 spill, spill 谁, 什么时候 split, 一个已经分配好的值是否应该被另一个值赶出寄存器.
+
+---
+
+## 3. 为什么不同虚拟寄存器可以共享同一个物理寄存器
+
+寄存器分配能够成立的关键, 在于一个函数中虽然可能出现上百甚至上千个虚拟寄存器, 但它们通常不会全部同时需要保存.
+
+考虑如下程序:
+
+```text
+v1 = a + b
+v2 = v1 * 2
+
+v3 = c + d
+v4 = v3 * 4
+```
+
+假设 `v1` 在第二条指令之后再也不会使用, 那么执行到:
+
+```text
+v3 = c + d
+```
+
+时, 保存 `v1` 的寄存器已经可以重新利用.
+
+于是完全可能得到:
+
+```text
+v1 → R0
+v3 → R0
+```
+
+这两个虚拟寄存器虽然名字不同, 却可以共享一个物理寄存器.
+
+决定它们能否共享寄存器的核心概念, 就是**活跃性**.
+
+一个值从定义开始, 到最后一次可能被使用的程序区域, 形成它的生命周期. 只要两个值在某个程序点同时需要保持有效, 它们就存在寄存器资源上的竞争关系.
+
+后面的第二部分会专门研究
+
+```text
+Liveness
+Live Range
+Live Interval
+```
+
+第三部分则进一步把这种竞争关系表示成
+
+```text
+Interference Graph
+```
+
+所以寄存器分配的知识链可以先建立成
+
+```text
+程序中的值
+    ↓
+什么时候仍然需要这个值
+    ↓
+Liveness
+    ↓
+哪些值会同时存活
+    ↓
+Interference
+    ↓
+哪些值不能共享寄存器
+    ↓
+Register Allocation
+```
+
+如果 "什么时候仍然需要一个值" 这个问题没有解决, 寄存器分配甚至无法判断两个虚拟寄存器能否安全地映射到同一个物理寄存器.
+
+---
+
+## 4. Register Pressure
+
+由此自然会出现另一个贯穿整个寄存器分配领域的概念: **register pressure**, 也就是寄存器压力.
+
+在一个程序点上, 如果有十个值同时需要保持有效, 那么可以直观地说, 这个位置产生了大约十个值的寄存器需求.
+
+如果机器在这一类值上只有八个可用寄存器:
+
+```text
+需要的 live values：10
+可用 registers：     8
+```
+
+这里就出现了明显的资源紧张.
+
+在现代编译器中实现还要考虑 register class, 固定寄存器, 别名关系以及某些寄存器暂时不可用等情况, 因此 register pressure 不能永远简单理解为"活跃变量数量". 然而这个直觉非常重要
+
+> 同时存活的值越多,  寄存器分配越困难。
+
+寄存器压力还揭示了一个重要事实: 寄存器分配从来都不完全独立于前面的代码生成过程.
+
+考虑:
+
+```text
+a = ...
+b = ...
+c = ...
+use a
+use b
+use c
+```
+
+如果指令调度把 `a`, `b`, `c` 的定义都提前, 那么三个值可能会长时间同时存活.
+
+另一种调度可能形成
+
+```text
+a = ...
+use a
+
+b = ...
+use b
+
+c = ...
+use c
+```
+
+这时峰值寄存器压力会明显降低.
+
+所以 Instruction Scheduling 会影响 Register Allocation; 反过来, 寄存器压力也会影响调度器应该怎样安排指令.
+
+这也是工业后端中一个非常重要的主题:
+
+```text
+Instruction Scheduling
+        ↕
+Register Pressure
+        ↕
+Register Allocation
+```
+
+到了 GPU 上, 这种关系还会更强. GPU 中每个线程占用的寄存器数量可能直接限制一个计算单元能够同时驻留多少个 warp 或 wave, 因此 register pressure 会进一步影响 occupancy. 这个问题我们会留到 GPU 寄存器分配部分完整展开.
+
+---
+
+## 5. 寄存器分配真正优化的是什么
+
+最容易形成的初始印象是:
+
+> 一个好的寄存器分配器应该尽量减少 spill.
+
+这个方向大体正确, 却还不足以描述真实的优化目标.
+
+首先, 不同 spill 的代价差异可能非常大.
+
+例如:
+
+```text
+for (...) {
+    use x
+}
+```
+
+如果 `x` 被 spill, 那么它可能在循环中产生大量动态 load.
+
+另一个只在冷路径上使用一次的值, 即使也发生 spill, 对程序运行时间的影响可能很小.
+
+于是编译器更关心的是:
+
+```text
+dynamic spill cost
+```
+
+也就是 spill 在实际执行中造成的成本.
+
+这也是为什么真实编译器经常使用 block frequency, loop depth, profile information 等信息来估计某个值有多值得占据寄存器.
+
+其次, 寄存器分配还要考虑 copy.
+
+假设前面的机器 IR 中存在:
+
+```text
+v2 = COPY v1
+```
+
+如果最终:
+
+```text
+v1 → R3
+v2 → R3
+```
+
+那么这条 copy 很可能完全消失.
+
+这种优化叫做 **coalescing**.
+
+于是一个 allocator 可能需要在两个选择之间权衡:
+
+```text
+减少 spill
+```
+
+以及:
+
+```text
+消除 copy
+```
+
+两者有时会互相冲突. 过度合并两个 live range 可能让它们形成一个更大的生命周期, 从而提高寄存器压力, 甚至导致新的 spill.
+
+继续往现实机器靠近, 还会出现更多目标.
+
+例如某个值跨越函数调用:
+
+```text
+v = ...
+call foo
+use v
+```
+
+如果把 `v` 放进一个 caller-saved register, 那么函数调用前后可能需要额外保存和恢复.
+
+如果把它安排到一个 callee-saved register, 又可能导致整个函数增加保存和恢复该物理寄存器的 prologue/epilogue 开销.
+
+因此一个寄存器选择本身也具有成本.
+
+最终, 一个工业寄存器分配器实际上面对的是多目标优化问题:
+
+```text
+spill cost
+copy cost
+register preference
+calling convention cost
+code size
+instruction constraints
+compile time
+```
+
+在 GPU 中还要进一步加入:
+
+```text
+occupancy
+```
+
+因此评价寄存器分配质量时, 只数 "spill 了多少个虚拟寄存器" 通常是不够的.
+
+---
+
+## 6. 寄存器分配在编译器后端中的位置
+
+从高度抽象的角度, 可以把代码生成流程理解成:
+
+```text
+LLVM IR / GIMPLE / other IR
+          ↓
+Instruction Selection
+          ↓
+Machine-level IR
+          ↓
+Register Allocation
+          ↓
+final machine instructions
+          ↓
+Machine Code
+```
+
+不过真正的后端流水线会复杂得多.
+
+以典型现代编译器为例, 指令选择完成后, 机器 IR 中经常已经出现目标体系结构的指令 opcode, 同时仍然保留虚拟寄存器.
+
+可以想象成:
+
+```text
+%0 = ADD64rr %1,   %2
+%3 = IMUL64rr %0,   %4
+```
+
+这里 `ADD64rr` 和 `IMUL64rr` 已经高度接近真实机器指令, 而 `%0`、`%1`、`%2`、`%3`、`%4` 仍然需要分配.
+
+寄存器分配完成之后,  它们可能变成:
+
+```text
+RAX = ADD64rr RBX,   RCX
+RDX = IMUL64rr RAX,   RSI
+```
+
+到这个阶段, 程序已经从 "拥有任意多虚拟寄存器的机器程序" 转化为 "受真实 ISA 资源限制的机器程序".
+
+Instruction Scheduling 和 RA 的相对顺序需要稍微谨慎处理.
+
+现代后端中经常同时存在:
+
+```text
+Pre-RA Scheduling
+Register Allocation
+Post-RA Scheduling
+```
+
+Pre-RA scheduling 仍然可以较自由地移动带虚拟寄存器的指令, 同时需要关注 register pressure.
+
+Post-RA scheduling 面对的已经是物理寄存器, 因此必须遵守更加具体的 hazard 和寄存器依赖.
+
+不同 target, 不同编译器甚至不同优化级别都可能采用不同 pipeline.
+
+不过这样概括地讲是没有问题的:
+
+> 寄存器分配通常工作在已经高度机器化的 IR 上, 它负责把虚拟寄存器世界收缩到真实机器允许的物理寄存器世界中.
+
+---
+
+## 7. 一个完整的小例子
+
+现在用一个很小的程序建立直觉.
+
+假设机器只有两个可用寄存器:
+
+```text
+R0
+R1
+```
+
+程序为:
+
+```text
+v1 = load a
+v2 = load b
+
+v3 = v1 + v2
+v4 = v3 * v1
+
+store v4
+```
+
+先观察 `v1`.
+
+它在第一条指令中被定义：
+
+```text
+v1 = load a
+```
+
+之后在:
+
+```text
+v3 = v1 + v2
+```
+
+中被使用, 同时在:
+
+```text
+v4 = v3 * v1
+```
+
+中还要再次使用.
+
+因此, 在计算完 `v3` 之后, `v1` 仍然需要继续保存.
+
+再看 `v2`.
+
+它只在:
+
+```text
+v3 = v1 + v2
+```
+
+中使用一次. 执行完这条指令以后, `v2` 就不再需要.
+
+于是可以产生一种分配:
+
+```text
+v1 → R0
+v2 → R1
+```
+
+执行:
+
+```text
+v3 = v1 + v2
+```
+
+以后, `v2` 已经死亡, 因此原来的 `R1` 可以拿给 `v3`:
+
+```text
+v3 → R1
+```
+
+此时:
+
+```text
+R0 = v1
+R1 = v3
+```
+
+正好满足下一条:
+
+```text
+v4 = v3 * v1
+```
+
+如果乘法允许结果覆盖其中一个已经死亡的 operand, 那么甚至可以继续复用寄存器.
+
+例如:
+
+```text
+v4 → R1
+```
+
+最后整个过程可能形成:
+
+```text
+v1 → R0
+v2 → R1
+v3 → R1
+v4 → R1
+```
+
+四个虚拟寄存器最终只需要两个物理寄存器.
+
+这里最重要的地方是资源复用发生的原因:
+
+```text
+v2 死亡
+   ↓
+R1 释放
+   ↓
+v3 使用 R1
+
+v3 死亡
+   ↓
+R1 再次释放
+   ↓
+v4 使用 R1
+```
+
+所以寄存器分配真正管理的是**随程序执行位置变化的资源占用**.
+
+---
+
+## 8. 再加一点压力
+
+现在稍微修改程序:
+
+```text
+v1 = load a
+v2 = load b
+
+v3 = v1 + v2
+v4 = v1 * v2
+v5 = v3 + v4
+
+store v5
+```
+
+还是只有:
+
+```text
+R0
+R1
+```
+
+两个物理寄存器.
+
+执行完:
+
+```text
+v3 = v1 + v2
+```
+
+以后, 事情开始变得麻烦.
+
+此时还需要:
+
+```text
+v1
+v2
+v3
+```
+
+因为下一条:
+
+```text
+v4 = v1 * v2
+```
+
+仍然要使用 `v1` 和 `v2`, 而最后的:
+
+```text
+v5 = v3 + v4
+```
+
+又要求 `v3` 继续保存.
+
+于是这里出现:
+
+```text
+3 个 simultaneously live values
+2 个 physical registers
+```
+
+寄存器资源已经无法同时容纳所有值.
+
+一个简单策略是 spill `v3`:
+
+```text
+v3 = v1 + v2
+store [slot],   v3
+
+v4 = v1 * v2
+
+load tmp,   [slot]
+v5 = tmp + v4
+```
+
+当然, 也可能选择 spill `v1` 或 `v2`.
+
+哪一个选择更好, 就开始涉及:
+
+```text
+谁以后还会被用多少次?
+位于什么控制流路径?
+是否处于循环内部?
+重新计算它是否比 load 更便宜?
+能否切断它的 live range?
+有没有别的寄存器可以换出来?
+```
+
+这正是寄存器分配算法真正开始发挥作用的位置.
+
+---
+
+## 9. Spill 也未必意味着整个值都长期待在内存里
+
+初学寄存器分配时还有一个很容易形成的简化模型:
+
+```text
+v 被 spill
+    ↓
+v 从此住在 stack
+```
+
+实际编译器通常会采取更加细粒度的处理.
+
+例如一个值的生命周期很长:
+
+```text
+v:
+
+|-------------------------------|
+```
+
+其中只有中间一小段寄存器压力特别大.
+
+那么编译器可能把它切成:
+
+```text
+|------|       |----------|
+ register        register
+        \_______/
+          spill
+```
+
+也就是说, 一个逻辑值的不同生命周期片段可以分别拥有不同存储位置.
+
+这就是后面非常重要的:
+
+> Live Range Splitting
+
+因此现代 RA 中的 "分配单位" 可能已经细化到 live range 的某一个 fragment, 而不再始终对应完整的源语言变量.
+
+这一点对理解 LLVM Greedy RA 尤其重要. LLVM Greedy 的很多能力都建立在 live interval splitting, eviction 和重新排队之上.
+
+---
+
+## 10. 源语言变量与寄存器分配对象并不是一一对应的
+
+再往前走一步, 我们还需要逐渐摆脱"一个 C 变量对应一个寄存器"的直觉. 有人经常说 C 语言就是高级汇编, 不是这样的.
+
+例如:
+
+```c
+int x = a + b;
+x = x * 2;
+x = x + c;
+```
+
+经过 SSA 化以后, 可能已经变成:
+
+```text
+x1 = a + b
+x2 = x1 * 2
+x3 = x2 + c
+```
+
+进入机器 IR 后, 由于 instruction selection, copy, two-address constraint, subregister, calling convention 等原因, 还可能出现更多机器级虚拟寄存器.
+
+于是寄存器分配器真正处理的对象通常更接近:
+
+```text
+machine-level values
+live ranges
+virtual registers
+```
+
+源语言变量与这些对象之间未必存在简单的一对一关系.
+
+同一个源变量可能对应多个虚拟寄存器; 同一个虚拟寄存器又可能在分配过程中被 split 成多个 live range fragment.
+
+因此以后讨论 RA 时, 我们会尽量使用:
+
+```text
+value
+virtual register
+live range
+live interval
+```
+
+这些机器后端层面的概念.
+
+---
+
+## 11. 为什么寄存器分配很难
+
+到目前为止, 我们已经可以看到 RA 的几个基本困难.
+
+首先, 物理寄存器数量有限, 同时存活的值却可能很多.
+
+其次, 值之间存在复杂的生命周期重叠, 因此不能任意共享寄存器.
+
+再次, 寄存器并不是全部等价. 一些指令只接受某类寄存器, 一些值要求 register pair, 一些机器存在 subregister alias, 一些指令隐式使用固定物理寄存器.
+
+此外, spill 的代价与执行频率相关, copy elimination 与 spill avoidance 之间可能发生冲突, calling convention 又会影响不同物理寄存器的使用成本.
+
+最终还需要考虑编译时间.
+
+一个非常昂贵的全局优化算法即使能够提高极少量代码质量, 也未必适合工业编译器.
+
+因此真实 RA 面对的是:
+
+```text
+有限寄存器资源
++
+复杂生命周期
++
+机器指令约束
++
+多种优化目标
++
+编译时间预算
+```
+
+后面我们讨论 Chaitin, Briggs, Linear Scan, LLVM Greedy, GCC IRA/LRA 时, 都可以把它们看成对这组约束的不同工程解法.
+
+---
+
+## 12. 第一部分需要建立的核心模型
+
+这一部分结束后, 最重要的是在脑中形成下面这个模型:
+
+```text
+前端 / 中端
+    ↓
+产生程序中的计算与数据依赖
+    ↓
+Instruction Selection
+    ↓
+Machine IR + Virtual Registers
+    ↓
+分析哪些值在什么时候仍然存活
+    ↓
+判断哪些值竞争同一个寄存器资源
+    ↓
+Register Allocation
+    ↓
+Physical Registers
+  ↙             ↘
+成功驻留         Spill / Split
+    ↓
+满足具体 ISA 约束
+    ↓
+Machine Code
+```
+
+从资源角度看, 它还可以压缩成:
+
+```text
+Virtual Registers
+        ↓
+    Liveness
+        ↓
+Register Pressure
+        ↓
+有限 Physical Registers
+        ↓
+┌───────────────┐
+│   Allocation  │
+└───────────────┘
+    ↓       ↓
+Register   Memory
+```
+
+之后所有复杂算法, 本质上都围绕几个问题展开:
+
+**哪些值会互相竞争?**
+
+**哪些值最值得拥有寄存器?**
+
+**拥有寄存器的值应该放在哪里?**
+
+**寄存器资源不足时应该牺牲谁?**
+
+**是否可以通过改变 live range, copy, 调度或者重新计算来缓解压力?**
+
+第一部分到这里完成了问题建模.
+
+第二部分开始, 我们就要进入整个寄存器分配理论真正的基础: **Liveness Analysis**。那里会正式回答一个最关键的问题——在程序的某一个位置上, 究竟怎样严格判断一个值 "还活着".
diff --git a/content/notes/register-allocation/gpu-occupancy/index.md b/content/notes/register-allocation/gpu-occupancy/index.md
new file mode 100644
index 0000000..ea000b7
--- /dev/null
+++ b/content/notes/register-allocation/gpu-occupancy/index.md
@@ -0,0 +1,384 @@
+---
+title: "第十部分: GPU 寄存器分配与 Occupancy"
+slug: "gpu-occupancy"
+lang: zh-Hans
+series: register-allocation
+weight: 110
+created: 2026-08-18
+updated: 2026-09-09
+license: CC-BY-SA-4.0
+---
+
+CPU 和 GPU 都需要把数量近乎无限的 virtual registers 映射到有限的 physical registers, 因而 liveness, interference, splitting, spilling 和 coalescing 这些基本问题并没有因为执行设备变成 GPU 而消失. 真正发生变化的是寄存器资源参与性能的方式. 由于 GCC 在 GPU 上已经落后很多了, 所以这一部分我们主要以 LLVM 来分析.
+
+CPU allocator 主要考虑单个 instruction stream 中的寄存器复用成本. GPU 同时维持大量 threads, 并把它们组织成 warp 或 wavefront 执行. 每个 thread 使用的寄存器最终都要占用 SM 或 CU 上有限的 register file. 当单个 kernel 的寄存器需求增加时, 一个执行单元能够同时驻留的 warps 或 waves 可能减少, 进而降低 GPU 隐藏 memory latency 和 execution latency 的能力. NVIDIA 的 CUDA 执行模型明确把 registers 和 shared memory 都视为决定 SM resident blocks 和 warps 数量的资源. AMD GPU 上的 SGPR 和 VGPR 同样从有限 register pools 中分配给正在执行的 wavefronts.
+
+所以 GPU register allocator 面对一个 CPU 上弱得多的反馈关系:
+
+```text
+register allocation
+        |
+        v
+register usage per thread / wave
+        |
+        v
+resident warps / waves
+        |
+        v
+occupancy
+        |
+        v
+latency hiding and throughput
+```
+
+这使 GPU RA 的优化目标发生了变化. 减少一次 spill 当然仍然有价值, 但如果为了避免 spill 而让整个 kernel 跨过一个 register-allocation threshold, 使 resident waves 从 4 个下降到 3 个, 最终性能可能反而下降. 反过来, 强行减少寄存器数量也可能引入大量 memory traffic 或额外 instructions. GPU RA 因此长期处在 register pressure, spill cost, instruction-level parallelism 和 occupancy 之间的权衡中. NVIDIA 的 `--maxrregcount` 文档直接把 per-thread register usage 和 thread parallelism 描述为一项 trade-off, AMD 的 GCN scheduler 也专门根据 register pressure 是否会降低 wave occupancy 调整调度策略.
+
+## 1 从 Thread Register 到整个 Register File
+
+先把 GPU 的执行层次和寄存器资源对应起来. CUDA 中的 threads 被组织成 thread blocks, block 进入 SM 后再被划分成 warps. 一个 warp 包含 32 个 threads, warp scheduler 从 ready warps 中选择下一条 instruction 发射. 每个 SM 拥有一组 32-bit registers, 这些寄存器资源会在驻留于该 SM 的 warps 之间分配.
+
+从程序员视角看, register 是 thread-private 的:
+
+```text
+thread 0: r0 r1 r2 ...
+thread 1: r0 r1 r2 ...
+thread 2: r0 r1 r2 ...
+...
+```
+
+但物理资源来自同一个有限 register file. 如果一个 warp 中每个 lane 都需要一个 32-bit value, 那么硬件需要为整个 warp 保存这一组 lane values. 所以 "kernel 使用 64 个 registers" 这类指标通常不是说整个 kernel 只占 64 个 hardware words, 而是在描述每个 thread 或每个 wave 的 architectural register requirement, 最终硬件会把它乘到同时驻留的执行上下文上. NVIDIA 也因此把 register usage 作为 per-thread kernel resource 报告, 并用它限制同时可以驻留的 threads 和 blocks.
+
+AMD 的 register organization 更能体现 GPU RA 与普通 CPU RA 的差别. AMD GCN/CDNA/RDNA 系列存在 SGPR 和 VGPR 两类主要 general-purpose register resources. VGPR 保存每个 lane 可以不同的 vector value, SGPR 保存编译器能够证明在整个 wave 中一致的 scalar value. 因而一个 wave-uniform 的值如果进入 SGPR, 整个 wave 只需要保存一份; 一个 divergent value 进入 VGPR 后, 每个 lane 拥有自己的 component.
+
+例如:
+
+```text
+wave:
+
+lane 0:   x = 3
+lane 1:   x = 7
+lane 2:   x = 5
+lane 3:   x = 9
+...
+```
+
+这种 `x` 是 divergent value, 需要 vector storage. 另一个值:
+
+```text
+wave:
+
+lane 0:   n = 128
+lane 1:   n = 128
+lane 2:   n = 128
+lane 3:   n = 128
+...
+```
+
+如果 compiler 能证明 `n` 在整个 wave 中 uniform, AMDGPU 可以让它走 scalar register bank. Uniformity analysis 因而和 register pressure 直接连接. 一个 value 从 VGPR 需求转化成 SGPR 需求时, 减少的不只是一个普通颜色冲突, 还改变了它消耗哪一类硬件资源. LLVM 的 AMDGPU register-bank logic 也明确区分 SGPR, VGPR 和 AGPR, 并要求只有已知 uniform 的 value 才能合法从 VGPR 语义转入 SGPR.
+
+部分 AMD GPU 还拥有 AGPR, 即 accumulator register resource, 用于 matrix accumulation 等操作. LLVM AMDGPU 的 pressure model 区分 SGPR, VGPR, AGPR 和 unified/mixed vector register categories, 并额外追踪 tuple pressure. 这意味着一个 AMDGPU kernel 的 register pressure 从一开始就不是一个单独整数.
+
+## 2 Register Pressure 为什么会改变 Occupancy
+
+Occupancy 表示一个 execution unit 上实际能够同时驻留多少 waves/warps 相对于硬件允许上限的程度. 限制 occupancy 的资源不只有 registers, 还包括 shared memory 或 LDS, thread-block/workgroup size, architectural wave slots 等. 可以把这种关系抽象写成 $W_{\text{resident}}=\min(W_{\text{arch}},W_{\text{reg}},W_{\text{smem}},W_{\text{block}},\ldots)$.
+
+其中 register-limited occupancy 可以粗略理解成 register file 容量除以每个 resident execution context 需要的 register storage. 假设某种抽象 GPU 每个 execution unit 有 $R$ 个可用 register units, 每个 wave 分配 $r$ 个 units, 那么单从寄存器容量得到的上限近似为 $\lfloor R/r\rfloor$. 真实硬件还会按照一定 granularity 分配 registers, 因而 occupancy 往往表现成阶梯函数, 而不是随着 register count 连续变化. LLVM AMDGPU 中专门存在 VGPR/SGPR allocation granule 以及根据 register 数量计算 waves-per-EU 的 target functions.
+
+用一个纯粹的抽象例子说明这个现象. 假设 register file 容量为 256 个 units, 分配粒度为 16. 如果一个 wave 需要 64 个 units, 最多可以容纳 4 个 waves. 如果最大 pressure 从 64 增长到 65, 硬件实际需要向上取整到 80, register-limited occupancy 就可能直接从 4 个 waves 下降到 3 个:
+
+```text
+logical pressure        allocated amount       resident waves
+
+      64                      64                     4
+      65                      80                     3
+      79                      80                     3
+      80                      80                     3
+```
+
+于是一个只增加了 1 个 live register 的代码变换, 有可能产生远大于 "多占一个寄存器" 的资源后果. 这类位置常被称为 occupancy cliff 或 occupancy threshold. Register allocator 和 scheduler真正关心的经常不是把 pressure 从 53 降到 52, 而是能否跨过一个会增加 resident-wave count 的 threshold.
+
+AMD 的公开架构资料能够直接看到这种效果. MI100 的 VGPR 和 SGPR 会动态分配给 wavefronts, VGPR 使用量直接限制并发 wavefront 数. AMD 给出的一个具体例子中, wave 使用 119 个 VGPR 时只能达到 2 个 active wavefronts.
+
+但 occupancy 不能单独当成 performance metric. 更多 resident waves 可以提供更多独立工作, 用一个 wave 的执行去覆盖另一个 wave 的 memory 或 pipeline latency. 当 kernel 本来已经有足够 latency hiding, 继续提高 occupancy 未必产生明显收益. 如果为了获得更高 occupancy 而减少 registers, compiler 还可能被迫增加 reload, rematerialization 或 dependency chains. NVIDIA 的 launch-bounds 文档甚至允许 compiler 在 occupancy requirement 已经满足后增加 register usage, 用更多 registers 换取更少 instructions 或更好的 single-thread latency hiding.
+
+所以 GPU allocator的目标不能写成简单的 "最小化 register count". 更接近实际问题的表达是: 在不引入过高 spill 和 instruction cost 的前提下, 控制 register usage 落在合适的 occupancy range 内.
+
+## 3 Liveness, Scheduling 与 Divergence
+
+GPU register pressure 仍然从 liveness 产生. 如果某个程序位置上有 80 个 VGPR-width values 同时 live, allocator 无法通过重新给这些 values 换 register number 来把 pressure 变成 40. 真正能够改变 maximum pressure 的办法仍然是缩短 live range, splitting, rematerialization, spilling, 改变 instruction schedule, 或者让某些 values 转移到其他 register bank.
+
+GPU instruction scheduling 因此和 RA 结合得非常紧. 假设有几个高 latency loads:
+
+```text
+load a
+load b
+load c
+
+compute a
+compute b
+compute c
+```
+
+把 loads 提前可以制造更多 memory-level parallelism, 让 GPU 同时等待多个 memory operations. 代价是 `a`, `b`, `c` 从 load 完成到真正消费之间都必须保持 live:
+
+```text
+a: |----------------|
+b:    |-------------|
+c:       |----------|
+```
+
+如果改成:
+
+```text
+load a
+compute a
+
+load b
+compute b
+
+load c
+compute c
+```
+
+live ranges 会缩短, register pressure 下降, 但 memory-level parallelism 和 instruction-level parallelism 也可能下降.
+
+AMDGPU 的 machine scheduler 直接把这个矛盾放进 scheduling policy. `GCNSchedStrategy` 会对一个 kernel进行多阶段 scheduling, 原因就是某些 scheduling regions 的 register pressure 会决定整个 kernel 的 occupancy. 当某个 region 接近会降低 occupancy 的 pressure threshold 时, scheduler 更倾向于压低 SGPR/VGPR pressure; 在不会影响 occupancy 的区域, 约束可以放松, 让 scheduler继续追求 ILP.
+
+这说明 GPU 上 "schedule first, register allocate later" 并不意味着两者彼此独立. Scheduler 虽然在 RA 前工作, 却必须预测 RA 后可能达到的 register usage. RA 随后又在 scheduler塑造出的 live ranges 上工作.
+
+Divergence 则进一步改变 register bank pressure. 一个基于 lane ID 的 value 通常会沿不同 lanes 产生不同结果, 后续计算也继续保持 divergent. 在 AMDGPU 上, 这种数据通常进入 VGPR. Workgroup ID, 某些 kernel arguments, constant addresses 或其他 wave-uniform data 则有机会进入 SGPR. AMD 的编程和 profiling 文档明确把 VGPR 描述为保存 wave 内不同 work-items 的数据, SGPR 保存编译期已知 wave-uniform 的数据.
+
+因此下面两种 IR 虽然从普通 SSA liveness 看都只有一个 value:
+
+```text
+v = thread_id.x + 1
+```
+
+和:
+
+```text
+v = workgroup_id.x + 1
+```
+
+对 AMD GPU 的 register allocation 可能完全不同. 前者随 lane 变化, 通常形成 vector value. 后者对一个 wave 中的 lanes 可以保持 uniform, 有机会使用 scalar resources. Uniformity analysis 做错或过于保守时, 本来只需要一个 scalar value 的数据可能长期占据 VGPR, 增加 vector pressure.
+
+SGPR 也不是无限资源. LLVM AMDGPU 的 occupancy model能够分别计算 SGPR 和 VGPR 对 occupancy 的影响, 某些架构代际上 SGPR allocation本身也会形成 occupancy limit. `GCNRegPressure` 因此比较 register-pressure states 时先比较 occupancy, 再比较 spilling, tuple pressure 和 raw register pressure, 而不是简单把所有 registers 加在一起.
+
+控制流 divergence 还会让 live range 形状更加复杂. 一个 value 可能只在部分 active lanes 上真正参与计算, 但 physical vector register仍然必须为整个 wave 保留对应 storage. AMD 的 `EXEC` mask 决定哪些 lanes 在当前 vector instruction 中 active, register allocation却不能把同一 VGPR 的不同 lanes随意分给无关 virtual values, 除非 backend 有专门的 lane/subregister packing 机制. 所以前面 CPU 上 "两个值在不同 CFG paths 上互斥, 可以共享寄存器" 的思想到了 SIMT 环境还需要同时理解 control-flow mask 和 register-lane semantics.
+
+## 4 LLVM AMDGPU 如何做 Register Allocation
+
+AMDGPU 的 physical register allocation发生在 LLVM machine-code backend 中. 上层可以经过 Clang, OpenMP/HIP lowering, MLIR GPU/ROCDL 等多种路径, 但进入 `lib/Target/AMDGPU` 后, allocator处理的已经是目标相关 Machine IR, 最终生成 AMD GPU ISA.
+
+AMDGPU 没有重新实现一套完全独立的 allocator. 它复用了 LLVM target-independent Greedy RA infrastructure, 同时把 GPU register banks 拆成多个 allocation stages. 优化编译中, SGPR 和 VGPR 分别运行 Greedy allocator, whole-wave/whole-quad mode 所需 registers 也有独立 allocation stage. AGPR 相关 values 在 RA 前经过专门 preparation, allocation 完成后还有 AGPR copy rewrite 等 target passes.
+
+主干可以简化成:
+
+```text
+Machine IR
+    |
+    v
+AMDGPU pre-RA optimizations
+and occupancy-aware scheduling
+    |
+    v
+prepare AGPR allocation
+    |
+    v
+SGPR register allocation
+    |
+    v
+lower SGPR spills
+    |
+    v
+WWM / special register allocation
+    |
+    v
+VGPR register allocation
+    |
+    v
+rewrite virtual registers
+    |
+    v
+post-RA AMDGPU passes
+```
+
+SGPR 和 VGPR 分开 allocation 很符合 AMD 硬件的资源结构. Scalar values 和 vector values 不只是两个不同 register classes, 它们拥有不同 storage semantics, 不同 spill possibilities, 对 occupancy 的作用也不同. Backend 因而可以在 SGPR allocation结束后先 lower 一部分 scalar spills, 再继续处理 VGPR allocation.
+
+Greedy RA 本身仍然使用第九部分讨论过的 `LiveInterval`, `LiveRegMatrix`, eviction 和 splitting. GPU-specific behavior主要通过 register classes, allocation order, target-specific spiller, scheduler 和 pressure model进入这套 framework. 也就是说, AMDGPU 并没有把 LLVM Greedy 换成一个名字叫 "GPU allocator" 的完全不同算法, 它改变的是 Greedy 所面对的 physical-resource topology 和 cost structure.
+
+`GCNRegPressure` 很适合观察这种 target-specific cost model. 它分别维护 SGPR, VGPR, AGPR 等 register kinds, 同时维护 tuple pressure. 比较两个 pressure states 时, 优先考虑哪个状态能够获得更高 occupancy, 随后才比较 spill risk, tuple pressure 和 raw register count.
+
+这里可以看到 CPU 和 GPU RA 的一个结构差异. CPU 上 allocator通常希望某个热点 region 使用较少 registers, 但一个函数最大 pressure 增加 1 不一定对全函数产生离散的执行资源变化. GPU 上一个 region 的 maximum VGPR pressure 可能决定整个 kernel 的 allocated register count, 最终改变所有 waves 的 residency. 所以 AMDGPU scheduler 和 allocator需要把局部 live-range decision 和 kernel-wide resource consequence联系起来.
+
+AGPR 又增加了一层 register-bank problem. 某些 matrix instructions 可以使用 accumulator registers, backend需要判断一些 values 保持在 VGPR 还是转换到 AGPR 更合适, 并考虑为这种转换插入 copies. AMDGPU 的 scheduling 和 RA machinery因此还会处理 VGPR/AGPR register-bank choice 和相关 copy cost.
+
+## 5 AMDGPU Spill 为什么比普通 Stack Spill 更复杂
+
+CPU spilling 常被简化成 "store 到 stack slot, use 前 load 回来". AMDGPU 同样有 scratch/private memory 作为寄存器溢出的 backing storage, 但不同 register banks 的 spill path 并不相同. AMDGPU backend为 scratch access准备相应的 private-segment state, VGPR spill 可以通过 scratch memory保存每个 lane 的 private value.
+
+VGPR spill 的直观形态可以画成:
+
+```text
+VGPR live range
+
+|-------------------------------|
+
+          high pressure
+               |
+               v
+
+      store lane value -> scratch
+              ...
+      load lane value <- scratch
+```
+
+Scratch 对每个 lane提供 private address semantics, 所以一个 spilled VGPR实际上对应 wave 中各 lanes 各自的 spilled value. Memory backing 和 register file 的组织方式不同, backend需要根据 wave size 和 scratch addressing convention产生正确地址.
+
+SGPR spill 更有 GPU 特征. 因为 SGPR 中保存的是整个 wave 共用的一份 scalar value, 直接把它按每个 lane复制到 scratch 会浪费 memory bandwidth. LLVM AMDGPU 可以优先把 SGPR spill 到 VGPR 的某个 lane, 使用 lane read/write operations 保存和恢复 scalar value. 如果最终必须进入 memory, backend可以先把 SGPR value写入 temporary VGPR lane, 再把 VGPR 保存到 scratch; restore 时反向执行.
+
+于是 SGPR spilling 可能形成这样的资源转换:
+
+```text
+SGPR pressure
+    |
+    | spill scalar value
+    v
+VGPR lane
+```
+
+如果 VGPR 也已经紧张, 才进一步变成:
+
+```text
+SGPR
+  |
+  v
+temporary VGPR lane
+  |
+  v
+scratch memory
+```
+
+这种设计减少 memory spill 的机会, 但它把 scalar-register pressure转移成 vector-register pressure. 假设 kernel 正好位于 VGPR occupancy threshold 附近, 多占一个 VGPR 可能比一次普通 SGPR spill 更昂贵. 所以 GPU spilling 的成本不能只按照 "load/store 次数" 建模, 还需要考虑 spill medium 本身占用了哪一类 register resource.
+
+AMDGPU 还存在 VGPR 与 AGPR 之间利用另一类 register storage 避免 memory spill 的机会. Backend 中有为 VGPR-to-AGPR 或 AGPR-to-VGPR spill保留寄存器的机制. 这使 register-bank pressure能够在一定条件下互相转移.
+
+Scratch 的使用本身也可能要求 kernel准备额外的 scratch-related SGPR state. AMDGPU backend 会根据 kernel 是否需要 private/scratch access建立相应的 scratch buffer 或 flat-scratch state. 因而 spilling 还可能反过来增加少量固定 register resource 和 prologue setup.
+
+这就是 GPU allocator里一个经常出现的反馈环:
+
+```text
+high VGPR pressure
+      |
+      v
+spill
+      |
+      v
+scratch instructions
+      |
+      v
+more address / temporary state
+      |
+      v
+new register pressure
+```
+
+好的 spiller需要避免让 spill rewrite 本身制造新的大规模 pressure problem. Rematerialization, splitting 和 spill-to-other-register-bank 都是在尝试切断这种反馈.
+
+## 6 NVIDIA: LLVM 到 PTX, 再到真正的 Physical Register Allocation
+
+NVIDIA 的 LLVM 路径与 AMDGPU 有一个架构上的根本差异. LLVM `NVPTX` backend 的目标是 PTX, 而 PTX 本身是 virtual ISA. LLVM NVPTX 把 LLVM IR 转换成 PTX assembly, PTX 中仍然可以声明大量 virtual registers. NVIDIA 的 PTX ISA甚至专门提供参数化 `.reg` 声明, 方便 compiler生成大量 `%r0`, `%r1`, `%r2` 这样的 virtual register names.
+
+例如 PTX 可以出现:
+
+```text
+.reg .b32 %r<100>;
+```
+
+这里声明的是 100 个 PTX virtual registers, 不是说 GPU hardware 已经为该 thread分配了 100 个 physical registers. PTX 之后还要经过 NVIDIA backend compiler.
+
+LLVM/NVIDIA pipeline 可以表示成:
+
+```text
+LLVM IR
+    |
+    v
+LLVM NVPTX backend
+    |
+    v
+PTX virtual ISA
+    |
+    +----------------------+
+    |                      |
+    v                      v
+  ptxas              CUDA Driver JIT
+    |                      |
+    +----------+-----------+
+               |
+               v
+            cubin
+               |
+               v
+       native GPU machine code
+```
+
+CUDA Driver 可以把 PTX JIT compile 成 native GPU machine code, `nvcc` 文档把 `ptxas` 定义为 PTX optimizing assembler. `ptxas --resource-usage` 的结果会报告最终使用多少 registers, 以及多少 spill loads/stores因为 variables 无法放进 physical registers 而产生. 因而 NVIDIA 路径中的最终 physical register assignment发生在 PTX optimizing backend这一侧, 不是 LLVM NVPTX 在生成 PTX 时完成.
+
+对 AMDGPU, LLVM Machine IR 中可以观察 Greedy allocator如何把 virtual VGPR/SGPR 映射成最终 architectural registers. 对 NVIDIA, LLVM NVPTX 输出仍然处于 virtual-register层, 最终 SASS register allocation由 NVIDIA backend完成. 修改 LLVM target-independent Greedy allocator不会直接替换 NVIDIA SASS allocator.
+
+NVIDIA 仍然把 register pressure控制接口暴露在 PTX 和 CUDA compilation model 中. PTX 的 `.maxnreg` 指令可以限制 kernel 每个 thread最多分配多少 registers, `.minnctapersm` 可以向 backend表达希望每个 SM 至少容纳多少 CTAs. 这些 directives 会让 optimizing backend在 per-thread register count 和 SM utilization之间做取舍.
+
+CUDA 层还有 `--maxrregcount` 和 `__launch_bounds__`. `--maxrregcount=N` 设置 GPU function 的 register upper bound. `__launch_bounds__(maxThreadsPerBlock,minBlocksPerMultiprocessor)` 则让 compiler根据希望同时驻留的 blocks数量推导 register budget. 如果原始 allocation 超过这个 budget, compiler会降低 register usage, 代价通常表现为更多 local-memory traffic 或更多 instructions.
+
+这提供了一个很直接的 RA experiment. 对同一个 kernel 编译多个版本:
+
+```text
+max registers/thread:
+
+128
+96
+80
+64
+48
+```
+
+然后观察:
+
+```text
+register count
+spill stores
+spill loads
+local memory usage
+occupancy
+kernel execution time
+```
+
+结果往往不会呈现 "register count 越低越快" 的单调关系. 在某个点之前, 减少 registers可能提高 occupancy而没有严重 spills; 继续压缩后, spill traffic开始主导性能. CUDA 文档也把 `maxrregcount` 明确描述为 individual-thread performance 和 thread parallelism 之间的 trade-off.
+
+NVIDIA register spill进入 local memory. CUDA 中的 local memory 是 thread-local address space, 但它的物理 backing 位于 device memory. Compiler在寄存器不足时可以把 value spill 到 local memory, 所以这种 spill拥有远高于 register access 的 latency, 同时经过 GPU memory hierarchy.
+
+PTX 还有 `.pred` 等不同 register types以及各种 special registers. 从 allocator角度看, 它们再次说明 GPU machine state 并不只是一个统一的整数 register file. 最终 NVIDIA backend还需要处理 predicate state, general registers, special architectural state以及不同 machine instructions 的 operand requirements. PTX virtual ISA把这些细节留给后端 compiler继续降低.
+
+## 7 CPU RA 与 GPU RA 的目标函数
+
+经过 AMDGPU 和 NVIDIA 两种实现以后, 可以把 GPU register allocation 和前面的 CPU register allocation放在同一个框架下比较.
+
+|问题|CPU|GPU|
+|---|---|---|
+|Register pressure 的主要后果|spill, copy, instruction constraints|spill之外还可能降低 occupancy|
+|一个 register assignment 的影响范围|主要影响当前 instruction stream|per-thread/per-wave usage 会乘到大量 resident execution contexts|
+|Spill backing|通常是 stack memory|private/local/scratch memory, 某些 GPU 还可利用其他 register bank|
+|Scheduler 与 RA|ILP, latency 和 pressure 权衡|ILP, latency, pressure 和 occupancy threshold 权衡|
+|Register classes|GPR, SIMD, special regs 等|scalar/vector/accumulator/predicate 等多类资源|
+|Control flow|CFG liveness|CFG liveness再叠加 SIMT divergence 和 execution mask|
+|评价指标|runtime, spills, copies, code size|runtime, spills, register count, occupancy, wave residency, code size|
+
+CPU allocator中一个常见目标是尽量避免 hot-path spill, 同时减少 copies 和 callee-saved cost. GPU allocator仍然追求这些目标, 但 register count本身开始成为 kernel-level resource metric. 一次看起来很便宜的 live-range extension可能让 maximum pressure跨过 occupancy threshold, 影响整个 kernel 的 resident-wave capacity.
+
+反过来, GPU allocator也不能把 occupancy最大化作为唯一目标. 一个 compute-heavy kernel如果拥有大量独立 arithmetic instructions, 更多 per-thread registers可以保存更多 intermediate values, 增加 ILP并避免 scratch traffic. 一个 memory-latency-bound kernel可能更愿意牺牲一些 single-thread efficiency, 换取更多 resident waves. AMD 的 occupancy-aware scheduler和 NVIDIA 的 launch-bounds机制都反映了这种取舍.
+
+所以评价 GPU register allocator时至少要把 register count和程序性能一起测量. 单独报告 "VGPR 从 72 降到 64" 这句话就没有包含足够信息. 如果 72 和 64 位于同一个 occupancy bucket, 这次 reduction可能没有任何 residency收益; 如果 64 刚好跨过一个 threshold, 效果可能很大. 同样, occupancy 从 50% 提高到 75% 也不能自动说明优化成功, 因为实现可能用大量 scratch spill换来了这个数字.
+
+GPU RA 的很多高级优化最终都围绕这种非线性成本展开. Scheduler可以重新排列 instructions 缩短 pressure peak, allocator可以围绕 high-pressure region splitting, rematerialization可以用 computation换 register residency, uniformity analysis可以把 VGPR pressure转移到 SGPR, register-bank choice可以在 VGPR 和 AGPR 之间重新分布资源. 有些研究甚至会有意识地接受少量 spill, 只为了把 maximum register count压过一个 occupancy threshold.
+
+到这里, 前十部分已经把 register allocation 从最简单的 virtual-to-physical mapping一路推进到了 CPU 和 GPU 的工业实现. 下一部分可以在这些基础上讨论更高层的协同优化问题: register-pressure-aware scheduling, profile-guided allocation, loop-aware splitting, rematerialization, instruction selection与 RA 的相互作用, occupancy-aware GPU allocation, allocator benchmark方法.
diff --git a/content/notes/register-allocation/graph-coloring/index.md b/content/notes/register-allocation/graph-coloring/index.md
new file mode 100644
index 0000000..02a7aa0
--- /dev/null
+++ b/content/notes/register-allocation/graph-coloring/index.md
@@ -0,0 +1,350 @@
+---
+title: "第四部分: 经典 Graph Coloring Register Allocation"
+slug: "graph-coloring"
+lang: zh-Hans
+series: register-allocation
+weight: 50
+created: 2026-08-14
+updated: 2026-09-09
+license: CC-BY-SA-4.0
+---
+
+上一部分已经把寄存器分配抽象成了图着色问题. 给定 interference graph $G=(V,E)$ 和 $K$ 个可用物理寄存器, allocator 希望给每个节点分配一种颜色, 并保证相邻节点颜色不同. 真正困难的地方在于, 一般图的 $K$-coloring 很难直接求解, 而编译器还需要同时决定哪些值值得 spill, spill 之后怎样改写程序, 以及如何在有限编译时间内获得足够好的结果.
+
+经典 graph coloring register allocation 的基本思路由 Chaitin 等人的工作奠定. 它利用一个简单的图论性质不断缩小 interference graph: 如果节点 $v$ 满足 $degree(v)<K$, 那么先暂时删除 $v$ 不会破坏剩余图的可着色性. 等剩余节点都处理完之后, 再按相反顺序恢复这些节点并选择颜色. 当图中已经找不到低 degree 节点时, allocator 需要选择一个可能被 spill 的节点来继续缩小图.
+
+这套流程通常可以概括为 Build, Simplify, Spill, Select 和 Rewrite. 不同文献对阶段名称和细节会有一些差异, 尤其是 Chaitin 原始方法和 Briggs 后来的 optimistic coloring 在 spill 的处理时机上并不完全相同. 先从最基本的 simplify 过程开始看.
+
+## 1 Simplify: 逐步拆掉容易着色的节点
+
+假设目标机器有 $K=3$ 个寄存器, 当前 interference graph 是:
+
+```text
+      a
+     / \
+    b---c
+    |
+    d
+```
+
+各节点 degree 为 $degree(a)=2$, $degree(b)=3$, $degree(c)=2$, $degree(d)=1$.
+
+因为 $d$ 满足 $degree(d)<3$, allocator 可以先把 $d$ 从图中移除, 同时把它压入一个栈中. 删除 $d$ 以后, $b$ 的 degree 从 3 降到 2. 这时图变成:
+
+```text
+      a
+     / \
+    b---c
+```
+
+现在 $a$, $b$, $c$ 的 degree 都是 2, 仍然小于 $K=3$. 可以继续删除其中任意一个, 例如先删 $a$, 再删 $b$, 最后删 $c$.
+
+删除顺序可能是:
+
+```text
+d
+a
+b
+c
+```
+
+如果每次删除都压入 stack, 最终得到:
+
+```text
+top
+ |
+ v
+c
+b
+a
+d
+```
+
+着色阶段会按相反顺序处理, 也就是先处理最后删除的节点.
+
+为什么这种删除是安全的, 第三部分已经给出了基本原因. 当删除节点 $v$ 时, 如果 $degree(v)<K$, 那么它最多有 $K-1$ 个邻居. 即使这些邻居最终使用了完全不同的颜色, 也最多占掉 $K-1$ 种颜色. 等 $v$ 被放回来时, 至少还有一种颜色可用.
+
+需要注意的是, simplify 时删除节点会不断改变其他节点的 degree. 一个最开始 degree 很高的节点, 可能随着邻居陆续删除而变成低 degree 节点. 因此算法关心的是当前简化图里的 degree, 不是最初 interference graph 中固定不变的 degree.
+
+例如:
+
+```text
+        a
+       /|\
+      b c d
+```
+
+如果 $K=3$, 开始时 $degree(a)=3$, 无法利用 $degree<K$ 的保证直接删除 $a$. 但如果 $b$ 和 $c$ 因为其他原因先被 simplify, $a$ 的当前 degree 就会下降到 1, 此时也可以安全移除.
+
+这也是 simplify 的主要价值. 算法没有直接解决整张图的 coloring, 而是利用局部容易处理的节点逐渐剥离图结构, 把真正困难的部分留到最后.
+
+## 2 当 Simplify 停下来时
+
+并不是所有图都能一直找到 $degree<K$ 的节点. 假设 $K=3$, interference graph 是完全图 $K_4$:
+
+```text
+      a------b
+      |\    /|
+      | \  / |
+      |  \/  |
+      |  /\  |
+      | /  \ |
+      |/    \|
+      c------d
+```
+
+每个节点的 degree 都是 3, 因此不存在 $degree<3$ 的节点. Simplify 在这里停住了.
+
+更一般地, 算法可能遇到一个 remaining graph, 其中所有节点都有 $degree\ge K$. 这并不能证明每个节点都必须 spill, 只能说明 simplify 所依赖的安全保证已经无法继续使用.
+
+Chaitin-style allocator 此时会选择一个节点作为 spill candidate, 暂时把它从图中移除. 删除这个节点以后, 它的邻居 degree 会下降, simplify 往往又可以继续进行.
+
+例如还是 $K_4$, 假设选择 $d$:
+
+```text
+      a
+     / \
+    b---c
+```
+
+剩余三个节点形成三角形, 在 $K=3$ 的情况下都满足 $degree=2<K$. 于是整张图又能继续 simplify.
+
+问题就转化成了: 应该选哪个节点作为 spill candidate?
+
+最简单的策略可以根据 degree 选择, 但实际效果通常不好. Spill 会在程序里产生额外的 load 和 store, 某个节点的动态使用频率往往比图结构本身更能决定 spill 的运行代价. 一个在冷路径上使用一次的值和一个位于内层循环中的值, 即使 degree 相同, spill 后的成本也可能相差几个数量级.
+
+因此经典 allocator 往往会估算一个 spill cost. 一个简单模型可以把 use 和 def 的执行频率累积起来, 再结合 degree 做归一化. 教材中经常看到类似 $\frac{spillCost(v)}{degree(v)}$ 的启发式指标. 如果这个值较小, 说明节点占据了较多冲突资源, 同时 spill 代价相对较低, 因而更适合作为候选.
+
+Spill cost 的具体公式没有唯一标准. 可以考虑 loop depth, block frequency, use 次数, def 次数, memory access cost 等因素. 工业实现还会利用 profile information 和 target-specific cost model. 图着色算法本身只要求在 simplify 卡住时能够选择一个节点继续推进, spill heuristic 决定的是这种选择对最终代码质量有多大影响.
+
+## 3 Chaitin 的 spill 与 Briggs 的 optimistic coloring
+
+这里需要区分两种处理 high-degree node 的方式.
+
+较早的 Chaitin-style 思路对 spill candidate 更悲观. 如果图中没有低 degree 节点, allocator 会选择一个节点准备 spill, 修改程序, 重新计算 liveness, 再重新构造 interference graph. 这种做法能够保证新的图更容易着色, 但也可能产生不必要的 spill.
+
+原因在第三部分已经见过. $degree(v)\ge K$ 只说明无法保证 $v$ 一定有颜色, 并没有证明最后真的没有颜色.
+
+Briggs 提出的 optimistic coloring 利用了这一点. 当 simplify 卡住时, allocator 仍然选出一个 potential spill candidate, 但先不立刻改写程序. 它把这个节点像普通节点一样从图中删除并压栈, 然后继续 simplify. 等到 select 阶段真正恢复这个节点时, 再检查它的邻居实际使用了哪些颜色.
+
+例如 $K=3$, 某个节点 $v$ 在 simplify 时有四个邻居:
+
+```text
+      a
+      |
+b ----v---- c
+      |
+      d
+```
+
+所以 $degree(v)=4\ge3$. 如果立刻根据 degree 判断 spill, 就会把 $v$ 写入内存.
+
+采用 optimistic coloring 后, $v$ 先作为 potential spill 被移除. 等 select 阶段重新处理 $v$ 时, 假设邻居最终分配为 $a\mapsto R_0$, $b\mapsto R_0$, $c\mapsto R_1$, $d\mapsto R_1$. 虽然 $v$ 有四个邻居, 它们实际只占用了两种颜色, 因此 $v$ 可以使用 $R_2$.
+
+只有当相邻节点真正覆盖了 $v$ 的全部可用颜色时, $v$ 才成为 actual spill. 比如 $K=3$, 最终相邻节点使用的颜色集合是 ${R_0,R_1,R_2}$, 那么 $v$ 就无法获得颜色.
+
+这也是 potential spill 和 actual spill 的区别. Potential spill 是 simplify 阶段为了继续缩小图而选择的高风险节点, actual spill 则是在 select 阶段确认无合法颜色之后产生的结果.
+
+Briggs 的这一改动看起来很小, 但它改变了算法看待 high-degree node 的方式. High degree 只参与启发式决策, 最后的 assignment 仍然由实际颜色占用决定.
+
+## 4 Select: 从栈中恢复节点并选择颜色
+
+Simplify 阶段结束时, interference graph 已经被逐步删除为空, 所有节点都保存在 stack 中. Select 阶段按照后进先出的顺序恢复节点.
+
+假设有 $K=3$, 颜色为 $R_0$, $R_1$, $R_2$, stack 是:
+
+```text
+top
+ |
+ v
+c
+b
+a
+d
+```
+
+首先弹出 $c$. 此时图中还没有已经着色的邻居, 可以给它选择任意合法颜色, 例如 $c\mapsto R_0$.
+
+随后弹出 $b$. 如果 $b$ 与 $c$ interfere, 那么 $R_0$ 已经被相邻节点占用, 可以选择 $b\mapsto R_1$.
+
+接着弹出 $a$. 如果 $a$ 同时和 $b$, $c$ 冲突, 那么它看到的已使用颜色是 ${R_0,R_1}$, 因此选择 $a\mapsto R_2$.
+
+最后恢复 $d$. 假设 $d$ 只和 $b$ 冲突, 那么它只需要避开 $R_1$, 可以使用 $R_0$ 或 $R_2$.
+
+这个过程可以写成简单的伪代码:
+
+```text
+while stack is not empty:
+    v = pop(stack)
+
+    forbidden = colors used by colored neighbors of v
+    available = legal_colors(v) - forbidden
+
+    if available is not empty:
+        color[v] = choose one color from available
+    else:
+        mark v as actual spill
+```
+
+在最简单的理论模型中, `legal_colors(v)` 就是全部 $K$ 个颜色. 有 register class, pre-colored registers 和 target constraints 以后, 每个节点的合法颜色集合可能不同.
+
+Select 也可以利用 register preference. 假设 $v$ 和某个节点之间存在 COPY relation, 并且对方已经获得 $R_1$, 那么在 $R_1$ 合法的情况下, allocator 可以优先选择 $R_1$, 从而消除 move. 这种偏好不会改变合法性约束, 但会影响最终生成的机器指令数量.
+
+在包含 optimistic coloring 的算法里, select 也是判断 potential spill 是否真正需要 spill 的阶段. 一个 high-degree node 被弹出时, allocator不再关心它最初的 degree, 只检查现在已经着色的邻居实际阻塞了哪些颜色.
+
+## 5 Rewrite: Spill 如何改变程序
+
+如果 select 阶段发现某个节点没有合法颜色, allocator 需要把它 spill. Spill 不能只在内部数据结构里给节点标一个 "memory" 颜色, 因为机器的大多数算术指令仍然要求 operand 位于寄存器中. Compiler 必须修改机器 IR, 在需要的位置插入 load 和 store.
+
+假设原来有:
+
+```text
+v = ...
+...
+x = v + 1
+...
+y = v * 2
+```
+
+如果 $v$ 被 spill, 可以给它分配一个 stack slot:
+
+```text
+spill_slot[v]
+```
+
+definition 之后把值写入内存:
+
+```text
+v_tmp = ...
+store v_tmp, [spill_slot]
+```
+
+每次 use 前再生成新的 temporary:
+
+```text
+load t1, [spill_slot]
+x = t1 + 1
+
+...
+
+load t2, [spill_slot]
+y = t2 * 2
+```
+
+原来 $v$ 可能拥有一个很长的 live range:
+
+```text
+v:
+|-----------------------------|
+```
+
+改写以后, 新生成的 $t1$, $t2$ 往往只有很短的生命周期:
+
+```text
+t1:
+       |---|
+
+t2:
+                       |---|
+```
+
+所以 rewrite 会直接改变 liveness 和 interference graph. 新 temporaries 需要重新分配寄存器, 但由于它们的 live ranges 较短, 通常比原来的长生命周期更容易着色.
+
+因此经典 allocator 在发生 actual spill 后通常需要重新开始一轮:
+
+```text
+Build
+  |
+  v
+Simplify
+  |
+  v
+Select
+  |
+  +---- no actual spill ----> finish
+  |
+ actual spill
+  |
+  v
+Rewrite
+  |
+  v
+recompute liveness
+  |
+  +-------------> Build again
+```
+
+这个迭代过程可能运行多轮. 如果第一次 rewrite 生成的新 temporaries 又导致新的冲突, 下一轮 allocation 可能继续选择其他 spill. 一个好的 spill heuristic 会尽量减少这种反复和最终产生的动态内存访问.
+
+Spill 也不一定必须真的执行 memory load. 如果一个值很容易重新计算, allocator 可以使用 rematerialization. 例如某个值只是常量 `0` 或一个便宜的地址计算, 在 use 位置重新生成它可能比从 stack load 更划算. 这部分会在后面的高级优化中展开.
+
+## 4.6 一个完整的着色过程
+
+下面用一个稍微复杂一点的图把 simplify, potential spill 和 select 串起来. 假设 $K=3$:
+
+```text
+       a
+      /|\
+     b-+-c
+      \|/
+       d
+       |
+       e
+```
+
+假设边集合为 $(a,b)$, $(a,c)$, $(a,d)$, $(b,c)$, $(b,d)$, $(c,d)$ 和 $(d,e)$. 其中 $a,b,c,d$ 构成 $K_4$, $e$ 只和 $d$ 冲突.
+
+初始 degree 为 $degree(a)=3$, $degree(b)=3$, $degree(c)=3$, $degree(d)=4$, $degree(e)=1$.
+
+因为 $degree(e)<3$, 先 simplify `e`. 删除以后 $degree(d)$ 从 4 降到 3. 剩余的 $a,b,c,d$ 每个 degree 都是 3, 没有节点满足 $degree<3$.
+
+此时 allocator 根据 spill heuristic 选择 $d$ 作为 potential spill, 将它移除并压栈. 删除 $d$ 以后, 剩余图成为:
+
+```text
+    a
+   / \
+  b---c
+```
+
+现在三个节点的 degree 都是 2, 因此可以继续 simplify. 假设依次删除 $a$, $b$, $c$. 整个删除顺序为:
+
+```text
+e
+d   <- potential spill
+a
+b
+c
+```
+
+stack 顶部是最后删除的 $c$:
+
+```text
+top
+ |
+ v
+c
+b
+a
+d
+e
+```
+
+Select 时先处理 $c$, 给它 $R_0$. 接着 $b$ 和 $c$ 冲突, 给 $b$ 分配 $R_1$. 然后恢复 $a$, 它同时和 $b$, $c$ 冲突, 因此获得 $R_2$.
+
+接下来轮到 potential spill 节点 $d$. 它和 $a$, $b$, $c$ 都冲突, 而这三个节点已经分别使用了 $R_2$, $R_1$, $R_0$. 此时 $d$ 的三个颜色全部被阻塞, 所以它确实成为 actual spill.
+
+最后恢复 $e$. 如果 $d$ 已经准备 spill, $e$ 的寄存器选择通常不会再受到 $d$ 的普通物理寄存器颜色限制, 因此可以正常获得寄存器.
+
+这个例子中的 $K_4$ 确实需要四种颜色, 所以某个节点最终无法在三个物理寄存器中完成 assignment. 如果换成另一个 high-degree graph, potential spill 节点的邻居可能复用颜色, 那么 select 时就可能避免 actual spill.
+
+这也说明 simplify 阶段和 select 阶段承担不同职责. Simplify 根据 degree 寻找结构上容易处理的节点, 在必要时做带风险的 potential spill 选择. Select 才看到邻居实际使用的颜色, 最终确定每个节点能否获得物理寄存器.
+
+## 7 Chaitin-Briggs 算法真正优化了什么
+
+纯粹的 graph K-coloring 只关心是否能找到合法颜色. Register allocator 还需要面对代码质量问题. 两个不同 coloring 都可能完全合法, 但其中一个产生更多 COPY, 另一个可能使用昂贵的 callee-saved register; 两种 spill 方案也可能产生完全不同的动态 load/store 数量.
+
+因此经典 graph-coloring allocator 通常有几个互相配合的启发式层面. Simplify 依据 degree 保持可着色性, spill selection 依据 cost model 选择比较合适的牺牲对象, select 阶段利用 register preference 决定具体物理寄存器, coalescing 则尝试消除 COPY. 这些机制共同作用, 最终获得的是一个可接受的近似解.
+
+这一点也能解释为什么 textbook Chaitin-Briggs 算法无法单独代表 GCC IRA 或 LLVM Greedy 这样的工业 allocator. Chaitin-Briggs 给出了 interference graph, degree, simplify 和 spill 的清晰理论框架, 工业实现还要处理 live range splitting, register classes, fixed constraints, rematerialization, eviction, scheduling interaction 等问题.
+
+下一部分会把目前暂时搁置的 COPY 问题完整展开. 在 interference graph 中, move-related nodes 希望获得相同颜色, aggressive coalescing 又可能让图更难着色. Briggs 和 George 提出的 conservative coalescing 条件, 以及 Iterated Register Coalescing 中的 Simplify, Coalesce, Freeze 和 SelectSpill, 会把 graph coloring allocator 进一步发展成一套完整的 move-aware 算法.
diff --git a/content/notes/register-allocation/interference-graph/index.md b/content/notes/register-allocation/interference-graph/index.md
new file mode 100644
index 0000000..a9762d5
--- /dev/null
+++ b/content/notes/register-allocation/interference-graph/index.md
@@ -0,0 +1,327 @@
+---
+title: "第三部分: Interference Graph 与冲突模型"
+slug: "interference-graph"
+lang: zh-Hans
+series: register-allocation
+weight: 40
+created: 2026-08-14
+updated: 2026-09-09
+license: CC-BY-SA-4.0
+---
+
+上一部分已经解决了 liveness 问题. 对任意一个程序位置, 我们可以判断哪些值仍然需要保留, 也可以进一步构造 live range 或 live interval. 寄存器分配接下来需要把这些位置相关的信息转化成寄存器之间的约束. 对经典 graph coloring allocator 来说, 这个约束通常通过 Interference Graph 表示.
+
+## 1 从 liveness 到 interference
+
+假设虚拟寄存器 $a$ 和 $b$ 在某个程序位置同时 live, 并且它们都需要使用同一类物理寄存器. 在这个位置上, 两个值必须同时保持有效, 因此 allocator 不能把它们分配到同一个物理寄存器. 我们称 $a$ 和 $b$ 之间存在 register interference.
+
+Interference Graph, 简称 IG, 就是把整个函数中的这种关系收集成一张无向图. 图中的节点代表 virtual register 或 allocator 实际处理的 live range, 边表示两个节点之间存在 interference. 例如有关系 $a-b$, $a-c$, $b-c$, $c-d$, 对应的图可以画成:
+
+```text
+      a
+     / \
+    b---c---d
+```
+
+假设目标机器有 $R_0$, $R_1$, $R_2$ 三个完全等价的物理寄存器. 一种合法分配是 $a\mapsto R_0$, $b\mapsto R_1$, $c\mapsto R_2$, $d\mapsto R_0$. 虽然 $a$ 和 $d$ 最终使用了同一个物理寄存器, 但图中没有 $a-d$ 这条边, 所以它们没有要求不同寄存器的生命周期重叠.
+
+从这个角度看, interference graph 是 liveness 的一种压缩. Liveness 保留了 "在哪个位置有哪些值同时 live" 这样的位置信息, IG 只记录 "哪两个值曾经发生过必须占用不同寄存器的冲突". 一旦图构建完成, allocator 可以暂时忽略冲突究竟发生在函数的哪一条指令附近, 直接研究节点之间的约束关系. 这很适合经典图着色算法, 但也意味着图本身会丢失冲突位置, 持续长度和执行频率等信息.
+
+如果某个程序位置的 live set 是 ${a,b,c}$, 那么三个值必须同时存在, 因而产生 $(a,b)$, $(a,c)$ 和 $(b,c)$ 三条边. 它们在图中形成一个大小为 3 的 clique:
+
+```text
+    a
+   / \
+  b---c
+```
+
+如果这三个值都只能使用同一个 register class, 那么这个位置至少需要三个物理寄存器才能让它们全部驻留在寄存器中. 因此某个位置的 simultaneously live values 会直接反映在 interference graph 的 clique 结构上.
+
+这里还需要把 "同时 live" 和 "竞争同一资源" 放在一起理解. 假设一个整数值只能使用 GPR, 一个浮点值只能使用完全独立的 FPR, 即使两者生命周期重叠, 它们也没有争夺同一个物理存储资源. 在最简单的理论模型里, 我们暂时假定所有节点属于同一个 register class, 所有物理寄存器也完全等价. 到真实机器约束部分再放宽这个假设.
+
+## 2 如何从 liveness 构造冲突图
+
+概念上最直接的方法是检查每个程序位置的 live set, 然后把其中的值两两连接. 如果某个位置有 $n$ 个 live values, 这种方法可能需要考虑 $O(n^2)$ 对关系. 经典实现通常利用 definition 来更直接地建立边.
+
+考虑一条普通三地址指令 $d=f(a,b)$. 假设这条指令执行完以后, $x$, $y$ 仍然 live, 同时新产生的 $d$ 也开始存活. 从这一时刻开始, $d$ 必须和 $x$, $y$ 共存, 因而需要建立 $(d,x)$ 和 $(d,y)$ 两条 interference edge.
+
+对普通指令可以概括成下面的关系:
+
+$$  
+d\in DEF(I),\ v\in LIVE\_OUT(I),\ d\neq v  
+\quad\Longrightarrow\quad  
+(d,v)\in E  
+$$
+
+这里 $E$ 是 interference graph 的边集. 这就是经常看到的 $DEF\times LIVE\_OUT$ 构图规则.
+
+为什么在 definition 的位置处理就足够捕获冲突, 可以从两个重叠的 live ranges 来理解. 假设 $a$ 已经处于 live 状态, 随后某条指令定义了 $b$, 并且 $a$ 在这条指令以后仍然需要保存. $b$ 从 definition 开始进入生命周期时, $a$ 还没有死亡, 因此两者的冲突会在 $b$ 的 definition 处暴露出来.
+
+例如:
+
+```text
+1: a = ...
+2: ...
+3: b = ...
+4: ...
+5: use a
+6: use b
+```
+
+执行完第 3 行以后, $a$ 仍然 live, $b$ 已经产生, 所以构图时会建立 $a-b$.
+
+现在手工走一遍稍完整的例子:
+
+```text
+1: a = load x
+2: b = load y
+3: c = a + b
+4: d = c * a
+5: e = d + b
+6: return e
+```
+
+根据上一部分的 liveness 分析, 各指令边界可以得到:
+
+| 位置  | Live before   | Live after    |
+| --- | ------------- | ------------- |
+| 1   | $\varnothing$ | ${a}$         |
+| 2   | ${a}$         | ${a,b}$       |
+| 3   | ${a,b}$       | ${a,b,c}$     |
+| 4   | ${a,b,c}$     | ${b,d}$       |
+| 5   | ${b,d}$       | ${e}$         |
+| 6   | ${e}$         | $\varnothing$ |
+
+第 1 行定义 $a$, 指令之后没有其他旧值需要和 $a$ 共存, 因此不增加边. 第 2 行定义 $b$, 此时 $a$ 仍然 live, 增加 $(a,b)$. 第 3 行定义 $c$, 指令之后 $a$ 和 $b$ 都仍然 live, 增加 $(c,a)$ 和 $(c,b)$. 第 4 行定义 $d$, 此时 $b$ 仍然 live, 增加 $(d,b)$. 第 5 行定义 $e$, 指令之后没有其他需要共存的虚拟寄存器.
+
+最终得到:
+
+```text
+      a
+     / \
+    b---c
+    |
+    d
+
+e
+```
+
+其中 $a$, $b$, $c$ 形成三角形. 如果目标机器只有两个可用寄存器, 这个子图无法完成 2-coloring. 如果有三个寄存器, 可以令 $a\mapsto R_0$, $b\mapsto R_1$, $c\mapsto R_2$, 而 $d$ 只和 $b$ 冲突, 所以可以使用 $R_0$ 或 $R_2$.
+
+从这个例子也可以看到一个细节. Interference edge 只表示两个值在某些位置必须使用不同资源, 它不记录冲突持续了多久. $a$ 和 $b$ 可能只重叠一条指令, 也可能在整个循环中长期重叠, 图上都只是同一条 $a-b$ 边. 以后决定 spill 谁时, allocator 还需要其他信息来补充这张图.
+
+## 3 Live range overlap, COPY 与 coalescing
+
+在简单模型中, 两个 live ranges 有实质性的重叠, 它们就会 interfere. 例如:
+
+```text
+a: |-------------|
+b:       |-------------|
+```
+
+这里 $a$ 和 $b$ 在中间一段同时 live, 如果竞争同一个 register class, 就需要不同物理寄存器. 如果生命周期完全错开:
+
+```text
+a: |-------|
+b:           |-------|
+```
+
+两者可以复用同一个寄存器.
+
+边界接触的情况需要更精确地看指令语义. 最典型的例子是:
+
+```text
+b = COPY a
+```
+
+假设这条 COPY 是 $a$ 的最后一次 use. 执行 COPY 时需要读取 $a$, COPY 之后旧的 $a$ 已经死亡, 新的 $b$ 开始继续保存这个值. 如果最后分配成 $a\mapsto R_0$, $b\mapsto R_0$, 那么这条指令会变成类似:
+
+```text
+R0 = COPY R0
+```
+
+随后可以删除.
+
+这说明构图时不能因为 source 和 destination 在 COPY 附近发生生命周期交接, 就机械地建立 $a-b$ interference edge. 一旦存在这条边, coloring 会强制 $color(a)\neq color(b)$, allocator 也就失去了通过寄存器分配消除这条 COPY 的机会.
+
+经典 graph-coloring allocator 对 COPY 通常会做特殊处理. 假设有 `b = COPY a`, 构造 $b$ 的冲突边时会把 source $a$ 从相关 live set 中排除, 同时记录 $a$ 和 $b$ 之间的 move relation:
+
+```text
+for COPY d <- s:
+    live = LIVE_OUT - {s}
+
+    for v in live:
+        addEdge(d, v)
+
+    recordMove(s, d)
+```
+
+于是图中同时存在两种不同性质的关系. Interference edge 表示两个节点不能使用同一个颜色, move relation 表示 allocator 希望两个节点尽量获得同一个颜色. 后者就是 register coalescing 的基础.
+
+看到 COPY 以后直接把两个节点合并也不安全. 假设 $Adj(a)={x,y}$, $Adj(b)={z,w}$. 如果把 $a$ 和 $b$ 合并为节点 $ab$, 新节点可能拥有 $Adj(ab)={x,y,z,w}$. 合并扩大了邻居集合, 节点的 degree 可能上升, 原本容易着色的图也可能因此变难.
+
+所以 coalescing 需要解决一个实际的权衡: 删除 move 可以减少机器指令, 合并 live ranges 又会改变 interference structure. 第五部分中的 Briggs criterion 和 George criterion, 都是在判断一次 coalescing 是否足够安全.
+
+COPY 还说明了为什么真实实现需要比 "每条指令一个离散位置" 更精细的位置模型. use 发生在什么时候, def 什么时候开始生效, two-address operand 或 early-clobber 会不会提前覆盖某个寄存器, 都可能影响两个生命周期是否真正需要同时占据不同资源. LLVM 的 SlotIndex 之类的数据结构就是为了表达这类细粒度位置关系.
+
+## 4 K-coloring 与 degree
+
+在理想模型里, 假设有 $K$ 个完全等价的物理寄存器, 每个 virtual register 恰好占一个寄存器, 所有 virtual registers 都可以使用这 $K$ 个寄存器. 令 interference graph 为 $G=(V,E)$, 颜色集合为 $C={R_0,R_1,\ldots,R_{K-1}}$. Register assignment 就是在寻找映射 $color:V\rightarrow C$, 满足下面的约束:
+
+$$  
+(u,v)\in E  
+\quad\Longrightarrow\quad  
+color(u)\neq color(v)  
+$$
+
+这就是 graph K-coloring 与寄存器分配之间的经典对应关系.
+
+一般图的 K-colorability 在 $K\ge 3$ 时是 NP-complete 问题. 对真实编译器来说, 函数中的 virtual registers 可能很多, allocator 还要同时考虑 spill cost, copy elimination, target constraints 和编译时间, 因此经典算法依赖启发式过程, 很少尝试直接求一般图的精确最优 coloring.
+
+在 Chaitin-style allocator 中, 节点 degree 是最常用的结构信息之一. $degree(v)=|Adj(v)|$, 表示 $v$ 和多少个其他节点发生 interference.
+
+如果 $degree(v)<K$, 那么在其他节点完成着色以后, $v$ 一定至少还剩一个可用颜色. 假设 $v$ 最多有 $K-1$ 个邻居. 即使这些邻居恰好分别使用了 $K-1$ 种不同颜色, 也只能占掉 $K-1$ 种颜色, 颜色集合中仍然至少剩下一种.
+
+例如 $K=4$, 且 $degree(v)=3$. 三个邻居最坏可以分别使用 $R_0$, $R_1$, $R_2$, 此时 $R_3$ 仍然可以分给 $v$. 这就是 simplify 操作能够成立的原因: allocator 可以暂时删除一个 $degree<K$ 的节点, 先处理剩余图, 最后再把这个节点放回来.
+
+当 $degree(v)\ge K$ 时, allocator失去了上述保证, 但节点仍然可能成功着色. Degree 统计邻居数量, 邻居最终可能复用颜色. 假设 $K=3$, $v$ 有三个邻居 $a$, $b$, $c$, 因此 $degree(v)=3$. 如果最后 $a\mapsto R_0$, $b\mapsto R_0$, $c\mapsto R_1$, 那么三个邻居实际只占用了两个颜色, $v$ 仍然可以使用 $R_2$.
+
+这也是 optimistic coloring 的基础. 当图中已经找不到 $degree<K$ 的节点时, allocator 可以选择一个 high-degree node 作为 potential spill candidate, 暂时将它移除并继续 simplify. 到 select 阶段恢复这个节点时, 再看相邻节点真正占用了多少种颜色. 如果相邻节点只使用了 $R_0$ 和 $R_1$, 节点仍然可以获得 $R_2$; 如果它的所有合法颜色都已经被占用, 此时才产生 actual spill.
+
+Potential spill 因此只是 simplify 阶段做出的一个风险选择. 它表示 allocator 暂时无法保证该节点以后一定能着色. 最后的 register assignment 仍然有机会成功.
+
+## 5 Clique, degree 与 register pressure
+
+Interference graph 和 register pressure 有直接联系, 但这两个概念描述的维度不同.
+
+如果某个程序位置同时有 $m$ 个 values live, 并且它们竞争相同寄存器资源, 这些节点会形成一个大小为 $m$ 的 clique. 如果机器只有 $K$ 个寄存器而某处出现了大小为 $K+1$ 的 clique, 这组 values 不可能全部同时获得寄存器.
+
+不过一个节点的 degree 不能直接理解成它所在位置的 register pressure. 假设 $v$ 生命周期很长, 在前半段与 $a$, $b$ 同时 live, 后半段与 $c$, $d$ 同时 live:
+
+```text
+early:
+{v, a, b}
+
+late:
+{v, c, d}
+```
+
+那么 $Adj(v)={a,b,c,d}$, 所以 $degree(v)=4$. 但程序中没有任何位置出现 ${v,a,b,c,d}$ 五个值同时 live. 这段程序的相关峰值 pressure 可能只有 3.
+
+Degree 是一个节点在整个生命周期中累计遇到的冲突数量. Register pressure 是某个具体程序位置上的瞬时资源需求. 长 live range 很容易拥有较高 degree, 因为它可能先后与很多互不重叠的短 live ranges 发生冲突.
+
+这也是仅按 degree 选择 spill candidate 容易出现问题的原因. 假设 $a$ 和 $b$ degree 相近, $a$ 只在冷路径中使用一次, $b$ 位于热循环中并被反复读取, spill 两者的动态成本显然不同. 经典 allocator 往往将使用频率和图结构结合起来, 例如考虑类似 $\frac{spillCost(v)}{degree(v)}$ 的指标. 工业实现还会继续加入 block frequency, loop depth, rematerialization cost 和 target-specific instruction cost 等信息.
+
+最大 clique 与图的 chromatic number 也不能简单等同. 一个大小为 $m$ 的 clique 至少需要 $m$ 种颜色, 因此 clique number 为 coloring 提供了下界. 对一般图来说, 即使没有大小为 $K+1$ 的 clique, 整张图仍然可能无法使用 $K$ 种颜色完成 coloring. 所以仅根据某一个位置的峰值 simultaneously live count, 无法完整解决一般 graph-coloring allocation.
+
+## 6 Spill 和 splitting 如何改变冲突结构
+
+如果当前 interference graph 很难着色, allocator 并不一定要一直在原图上寻找更复杂的颜色组合. Spill 和 live range splitting 都会改变程序的生命周期结构, 从而产生一张不同的冲突图.
+
+假设原来有一个很长的 live range:
+
+```text
+v:
+|------------------------------|
+```
+
+如果把整个 $v$ spill 到 stack, 原来对 $v$ 的 use 可能被改写成:
+
+```text
+load t1, [slot]
+use t1
+
+...
+
+load t2, [slot]
+use t2
+```
+
+原来的长生命周期被若干短 temporary 取代:
+
+```text
+t1:
+|---|
+
+t2:
+                    |---|
+```
+
+`t1` 只和第一个局部区域里的 values 冲突, `t2` 只和第二个区域里的 values 冲突. 因为 rewrite 已经改变 def/use 和 liveness, 原来的 interference graph 也需要重新计算. 传统 graph-coloring allocator 因此经常形成这样的迭代过程:
+
+```text
+Build graph
+    |
+    v
+Simplify / Color
+    |
+    v
+Choose spill
+    |
+    v
+Rewrite program
+    |
+    v
+Recompute liveness
+    |
+    +------> Build graph again
+```
+
+Live range splitting 做得更细. 假设 $v$ 前半段与 $a,b,c$ 冲突, 中间与 $d,e,f$ 冲突, 后半段与 $g,h$ 冲突. 如果始终把 $v$ 当作一个节点, 它会有 $Adj(v)={a,b,c,d,e,f,g,h}$. 如果把生命周期拆成三个 fragments:
+
+```text
+v1:
+|---------|
+
+v2:
+           |----------|
+
+v3:
+                       |--------|
+```
+
+那么可以得到更局部的冲突集合, 例如 $Adj(v_1)={a,b,c}$, $Adj(v_2)={d,e,f}$, $Adj(v_3)={g,h}$. 三个 fragments 可以分别做 allocation decision, 甚至分别使用不同物理寄存器, 中间通过 copy 或 memory location 连接.
+
+这类操作说明 register allocation 的输入结构本身是可以被修改的. 现代 allocator 经常通过 splitting, rematerialization, spill rewrite 等方式重新塑造 live ranges, 再解决一个更容易的局部 allocation problem.
+
+## 7 真实机器对简单 coloring 模型的扩展
+
+前面的 K-coloring 模型假设每个节点拥有同一组颜色. Register class 一出现, 这个假设就开始放宽. 假设机器有 $GPR={R_0,R_1,R_2,R_3}$ 和 $FPR={F_0,F_1,F_2,F_3}$, 整数 virtual register 只能选择 GPR, 浮点 virtual register 只能选择 FPR. 不同节点因此会拥有不同的 candidate set.
+
+还可能出现部分重叠的候选集合. 例如 $Candidates(a)={R_0,R_1,R_2}$, $Candidates(b)={R_1,R_2,R_3}$. 这时 allocator 除了处理 $a$ 与 $b$ 是否 interfere, 还需要处理各自允许选择哪些颜色.
+
+Pre-colored register 又引入固定颜色. 某些 ABI 或指令约束会要求一个值位于指定物理寄存器. 如果节点 $p$ 已经固定为 $color(p)=RAX$, 任何与 $p$ 冲突的节点都必须避开 $RAX$. 如果有:
+
+```text
+RAX = COPY v
+```
+
+allocator 反而会希望在合法的情况下直接令 $v\mapsto RAX$, 从而消除这条 COPY. 所以工业 assignment 同时包含硬约束和 register preference.
+
+Register aliasing 会让颜色本身发生重叠. x86 的 `RAX`, `EAX`, `AX`, `AL` 等寄存器共享底层硬件状态. 即使它们名字不同, 也不能被当成四个完全独立的颜色. 某些 value 还会占用 register pair 或 register tuple, 例如一个值可能要求 $(R_2,R_3)$ 两个连续寄存器, GPU vector value 也可能要求连续 VGPR tuple. 连续性, 对齐和 subregister constraint 都需要额外的 target-specific 表示.
+
+因此 textbook K-coloring 提供的是 register interference 的基本模型. 当机器约束逐渐加入以后, 实际问题会演化成带有不同 candidate sets, fixed assignments, aliasing, tuples 和 assignment costs 的受约束资源分配问题.
+
+## 8 显式冲突图与按需 interference query
+
+经典 Chaitin-Briggs allocator 通常会显式保存 interference graph. 对每个节点维护 $Adj(v)$ 和 $degree(v)$, simplify 时删除节点并更新邻居 degree, select 时查看已经着色的邻居使用了哪些颜色. 这种结构和算法本身非常契合.
+
+显式图也有成本. 如果函数中有 $N$ 个 virtual registers, interference edge 在最坏情况下可以接近 $O(N^2)$. 而且图只告诉 allocator 两个值曾经冲突过, 不能直接告诉它们在哪段 live range 上发生重叠.
+
+使用 live intervals 的 allocator 可以采用另一种方式. 假设:
+
+```text
+a:
+[2, 10)      [20, 30)
+
+b:
+      [8, 15)
+```
+
+因为区间 $[2,10)$ 和 $[8,15)$ 有交集, allocator 可以在真正需要检查时发现 $a$ 与 $b$ interfere. 这种方式把时间位置信息保留下来, 对 splitting 和局部冲突查询更方便, 同时避免预先物化整张全局图.
+
+LLVM Greedy Register Allocator 大量使用 LiveIntervals 和按需 interference checking, 整体结构和 textbook Chaitin allocator 有明显差异. GCC IRA 保留了更直接的 conflict 和 coloring 体系, 但也围绕 allocno, object, region 等抽象扩展了经典图模型.
+
+因此后面看到 "interference" 时需要区分两个层次. 一层是数学上的关系: 两个 live ranges 不能同时占用同一份物理寄存器资源. 另一层是实现选择: compiler 可以把这种关系显式存成 graph edge, 也可以利用 live range 数据结构在需要时计算.
+
+下一部分将沿着显式 interference graph 继续, 进入 Chaitin 和 Briggs 的经典 Graph Coloring Register Allocation. Build 阶段已经在这一部分基本解决, 接下来要研究 simplify 如何依靠 $degree<K$ 逐步缩小图, 图无法继续 simplify 时如何选择 potential spill, 以及 select 阶段为什么仍然可能把 high-degree 节点成功着色.
diff --git a/content/notes/register-allocation/linear-scan/index.md b/content/notes/register-allocation/linear-scan/index.md
new file mode 100644
index 0000000..6eb82c0
--- /dev/null
+++ b/content/notes/register-allocation/linear-scan/index.md
@@ -0,0 +1,278 @@
+---
+title: "第六部分: Linear Scan Register Allocation"
+slug: "linear-scan"
+lang: zh-Hans
+series: register-allocation
+weight: 70
+created: 2026-08-17
+updated: 2026-09-09
+license: CC-BY-SA-4.0
+---
+
+前面两部分一直沿着 interference graph 展开. Graph coloring 把 liveness 压缩成 "哪些 live ranges 互相冲突", 然后在这张图上寻找颜色. Linear Scan 选择保留另一种信息: 每个值在程序的哪些位置 live. 它先把机器指令排列在线性的 position space 中, 再把每个 value 的生命周期表示成 live interval, 按生命周期出现的顺序向前扫描.
+
+Linear Scan 最初受到重视, 很大程度上来自它的编译速度和实现复杂度. 它尤其适合 JIT 场景, 因为 JIT 编译器需要把编译延迟控制得很低. 但旦 "Linear Scan" 并不意味着所有现代实现都只是一个极简单的线性循环. 一加入 lifetime holes, interval splitting, register constraints, fixed intervals 和更复杂的 spill heuristic, 算法也会逐渐形成相当丰富的数据结构. LLVM 自己曾长期使用 Linear Scan 作为默认优化寄存器分配器, 从 2004 年一直使用到 LLVM 3.0 前后; LLVM 3.0 切换到新的 Greedy allocator 时, 官方回顾也明确指出 Linear Scan 在 LLVM 中已经工作了很多年, 同时暴露出了全局 splitting 和 allocation order 上的限制.
+
+## 1 从 live interval 开始
+
+假设机器指令已经按照某种线性顺序编号:
+
+```text
+0    4    8    12    16    20    24    28
+|----|----|-----|-----|-----|-----|-----|
+```
+
+几个 virtual registers 的生命周期可能是:
+
+```text
+v1:  |----------------|
+v2:       |------|
+v3:                    |---------|
+v4:            |-----------------|
+```
+
+为了先讲清经典算法, 暂时把每个 value 看成一个连续区间 $[start,end)$. Linear Scan 首先按照 $start$ 从小到大排列所有 intervals. 扫描到一个新 interval 时, allocator 只需要知道哪些旧 intervals 在当前位置仍然 live. 这些 intervals 构成 `active` 集合.
+
+例如当前处理 $v_3$, 它从位置 20 开始. 如果 $v_1$ 在位置 16 已经结束, 那么 $v_1$ 占用的物理寄存器可以立即释放. 如果 $v_4$ 一直活到位置 28, 它仍然留在 `active` 中. 因此 Linear Scan 不需要预先建立一张完整的 interference graph. 两个 intervals 是否竞争寄存器, 可以根据它们在扫描当前位置是否同时 active 来判断.
+
+经典算法的骨架大致如下:
+
+```text
+sort intervals by increasing start position
+
+for each interval i:
+    ExpireOldIntervals(i)
+
+    if active contains fewer than K intervals:
+        assign a free register to i
+        add i to active
+    else:
+        SpillAtInterval(i)
+```
+
+这里 `active` 通常按照 interval 的结束位置排序. 这样 `ExpireOldIntervals` 可以从最早结束的 interval 开始检查. 如果某个 active interval $j$ 满足 $end(j)\le start(i)$, 那么 $j$ 已经在新 interval 开始之前死亡, 它的寄存器可以回收到 free register 集合中.
+
+假设机器只有三个寄存器 $R_0,R_1,R_2$, intervals 是:
+
+```text
+v1: [0,  8)
+v2: [2, 12)
+v3: [4, 10)
+v4: [9, 16)
+```
+
+处理 $v_1$ 时, `active` 为空, 可以令 $v_1\mapsto R_0$. 到 $v_2$ 时, $v_1$ 还没结束, 所以令 $v_2\mapsto R_1$. $v_3$ 开始于 4, 此时 $v_1$, $v_2$ 都 active, 因而使用最后一个空闲寄存器 $R_2$.
+
+扫描到 $v_4$ 时, $start(v_4)=9$. $v_1$ 已经在位置 8 结束, `ExpireOldIntervals` 会把它移出 active 并释放 $R_0$. 因此 $v_4$ 可以直接使用 $R_0$. 整个过程只需要维护当前仍然覆盖扫描位置的 intervals.
+
+这种视角和 interference graph 有一个明显区别. Graph coloring 会提前记录 $v_i$ 与哪些其他 values 冲突; Linear Scan 更关心当前扫描位置上哪些 intervals 还没有结束. 对 straight-line interval model 来说, temporal overlap 本身就足够告诉 allocator 当前有哪些寄存器被占据.
+
+## 2 当所有寄存器都被占用
+
+真正需要做决策的情况发生在 `active` 已经包含 $K$ 个 intervals, 又有一个新 interval 开始时.
+
+假设只有三个寄存器:
+
+```text
+v1: |------------------------|   end = 30
+v2:    |----------|              end = 16
+v3:       |------------|         end = 22
+v4:          |-----|             end = 14
+             ^
+             scan reaches v4
+```
+
+当 $v_4$ 开始时, $v_1$, $v_2$, $v_3$ 都还 active. 四个 intervals 在当前位置同时需要资源, 三个寄存器无法全部容纳. 最简单的 Linear Scan spill heuristic 会比较当前 interval 和 active intervals 的结束位置. 如果 active 中结束最晚的 interval 比当前 interval 活得更久, 可以 spill 那个长 interval, 把它的寄存器让给当前 interval; 如果当前 interval 自己结束得最晚, 则直接 spill 当前 interval.
+
+在这个例子中, $v_1$ 的 $end=30$, 而 $v_4$ 的 $end=14$, 因此一种合理选择是暂时牺牲 $v_1$. 这样 $v_4$ 很快就会死亡, 寄存器随后又可以重新使用. 如果反过来把 $v_4$ spill, 只是为了让 $v_1$ 从当前位置一直占据寄存器到 30, 可能造成更多后续冲突.
+
+经典伪代码经常写成类似:
+
+```text
+spill = active interval with latest end
+
+if end(spill) > end(current):
+    current gets spill's register
+    spill goes to memory
+    remove spill from active
+    add current to active
+else:
+    current goes to memory
+```
+
+"结束得最晚" 是一个很局部的 heuristic. 它利用了扫描顺序提供的信息: 一个很晚才结束的 interval 会长期占据资源, 把它移走可以更快地释放寄存器容量. 更成熟的实现会考虑 use position, spill cost 和 split point. 一个 interval 虽然结束得很晚, 但可能在当前点之后马上就要使用; 另一个 interval虽然整体较短, 下一次 use 却还很远. 只看最终的 `end` 会错过这种区别.
+
+这也是 Linear Scan 后来逐渐发展出 lifetime position based spill decisions 的原因. allocator 可以问 "这个 interval 下一次什么时候真正需要寄存器", 而不只是 "它最后什么时候结束". 如果一个 value 在很长一段距离内都没有 use, 当前寄存器紧张时先把它移出去通常很自然.
+
+## 3 Lifetime holes 和 inactive set
+
+第二部分已经提过, 真实 live range 往往不是单一连续区间. 一个 live interval 可以由多个 segments 组成:
+
+```text
+v1:
+|---------|          |-----------|
+
+v2:
+           |--------|
+```
+
+如果只取 $v_1$ 最早的 start 和最晚的 end, 会得到:
+
+```text
+v1:
+|--------------------------------|
+```
+
+这样会把中间的 lifetime hole 也当成寄存器占用区间, 从而制造假的 interference.
+
+支持 holes 的 Linear Scan 通常除了 `active` 之外再维护一个 `inactive` 集合. 当一个 interval 的整个生命周期尚未结束, 但当前位置恰好处于它的 hole 中时, 它从 `active` 移到 `inactive`. 当扫描位置再次进入它的下一个 live segment 时, interval 又可能从 `inactive` 回到 `active`.
+
+例如:
+
+```text
+v1: |--------|          |--------|
+v2:      |----------------|
+              ^
+              current position
+```
+
+在箭头位置, $v_1$ 整个 interval 还没有结束, 因为右边还有第二个 segment. 但它当前并不 live, 所以不应该继续占据 active 资源. $v_2$ 在这一段可以暂时使用某个不会与 $v_1$ 后续 segment 发生冲突的寄存器; 如果两者后面的 segments 再次重叠, allocator仍然需要提前考虑这种未来冲突.
+
+因此支持 holes 以后, `inactive` 不能被简单视为 "完全不相关". 一个 inactive interval 当前不占据寄存器, 但它未来可能重新变为 active. 当 allocator 为新 interval 选择 physical register 时, 还要检查使用这个寄存器的 inactive intervals 会不会在未来与当前 interval 相交. 这会影响某个物理寄存器 "能够安全使用到哪个位置".
+
+现代 Linear Scan 常用 `free_pos[reg]` 一类信息描述这种情况. 如果一个寄存器当前完全空闲, 但未来在位置 40 会与某个 inactive interval 冲突, 那么这个寄存器对当前 interval 可以自由使用到位置 40. 如果当前 interval 会一直活到 80, allocator 可以在 40 附近 split 它, 让前半段使用这个寄存器, 再处理后半段.
+
+至此, Linear Scan 已经开始从 "给整个 interval 一次性分配一个寄存器" 发展成对 interval fragments 做决策.
+
+## 4 Live range splitting
+
+早期 Linear Scan 的主要弱点之一, 就是一个长 interval 在寄存器不足时很容易被整体 spill. 假设:
+
+```text
+v:
+|--------------------------------------|
+
+          high pressure
+              |
+              v
+         |----------|
+```
+
+如果只有中间一小段发生寄存器冲突, 将整个 $v$ 从 definition 到最后 use 都 spill 到 memory 会浪费大量本来可以驻留寄存器的区域.
+
+Splitting 可以把它改成:
+
+```text
+v1:
+|-------------|
+
+v2:
+              |----------|
+
+v3:
+                         |-------------|
+```
+
+然后分别决定 $v_1$, $v_2$, $v_3$ 的 location. 例如两边驻留寄存器, 中间高压力区域进入 stack:
+
+```text
+register        stack          register
+|-------------|----------|----------------|
+```
+
+这样真正执行的 spill/fill 只发生在分界附近.
+
+Linear Scan 中的关键问题变成 split point 放在哪里. 如果当前 physical register 只能安全使用到位置 $p$, 可以把 interval 在 $p$ 附近切开. 如果某个 child interval 下一次 use 很晚, 也可以把它推迟到靠近 use 的位置才 reload. 更高级的算法会综合 block boundaries, loops, calls 和 use positions 选择 split point.
+
+这一步显著缩小了简单 Linear Scan 和高质量 allocator 之间的差距. 很多时候, allocator 质量差异并不来自 "颜色怎么选", 而来自 live range 被怎样拆开以及 spill/fill 被放在哪里.
+
+LLVM 从旧 Linear Scan 切换到 LLVM 3.0 Greedy allocator 时, 官方总结恰好提供了一个很典型的工业案例. LLVM 的旧 Linear Scan 按线性顺序访问 live ranges, 使用 active list 检查 interference. 一旦所有寄存器被 active intervals 阻塞, 它需要选择 live range spill. LLVM 当时的实现很难做完整的 global live range splitting, 因为新产生的 fragments 如果应该被放回已经扫描过的程序区域, 就需要回退 allocation 状态. LLVM 开发者认为这种 backtracking 与 Linear Scan 的整体结构很不协调.
+
+## 5 Fixed intervals 和机器寄存器约束
+
+到目前为止, 我们假设所有 physical registers 都只被 virtual intervals 占用. 真实机器还会有固定寄存器使用.
+
+假设某条指令在位置 24 必须使用 `RAX`, 那么可以把这个约束看作 `RAX` 自己存在一个 fixed interval:
+
+```text
+RAX fixed:
+                    |--|
+
+v1:
+          |------------------|
+```
+
+如果准备把 $v_1$ 分配到 `RAX`, 两者在位置 24 会冲突. allocator 可以选择其他 physical register, 也可以在固定使用之前 split $v_1$, 让它避开 `RAX` 被占据的区域.
+
+函数调用也可以用类似方式理解. 一次 call 会 clobber 一组 caller-saved registers. 对跨 call 的 live interval 来说, 这些寄存器在 call 位置形成固定障碍:
+
+```text
+v:
+|----------------------------|
+
+              CALL
+               |
+               v
+
+R0:           |x|
+R1:           |x|
+R2:           |x|
+```
+
+如果 $v$ 需要跨 call 保持, allocator 可以选择未被 call clobber 的寄存器, 也可以在 call 周围 split, spill/fill, 或根据具体成本做其他安排.
+
+Two-address instruction, tied operand 和 pre-colored value 最终也可以转化成 interval assignment 上的约束. 因此成熟的 Linear Scan 同样需要真正理解目标 ISA, 并不因为它没有显式 interference graph 就能绕过 machine constraints.
+
+## 6 Linear Scan 为什么快
+
+Linear Scan 的基本优势来自它处理 interference 的方式. Graph-coloring allocator 往往需要显式或逻辑上处理全局冲突关系, 而经典 Linear Scan 只维护扫描点附近的 active intervals. 如果 intervals 已经按 start position 排好序, 每个 interval 只需要被插入和移出少数几个工作集合, physical register 的数量通常又远小于 virtual register 数量, 因而整个核心过程可以非常高效.
+
+"Linear" 这个名字需要稍微宽松地理解. intervals 的排序可能需要 $O(n\log n)$; active set 如果用平衡结构维护也有额外成本; splitting 会产生新的 intervals; 带 lifetime holes 的实现还要管理 inactive sets 和 future intersection queries. 因此一个工业 Linear Scan allocator 的实际复杂度不能简单写成严格的 $O(n)$ 然后结束讨论. 它的优势更多来自算法只维护相对局部和有序的状态, 避免构造和反复修改庞大的全局 interference graph.
+
+这也是它长期受到 JIT 编译器欢迎的原因. JIT 对编译延迟高度敏感, 一个能够快速产生足够好代码的 allocator 往往比追求更复杂全局优化更合适. AOT 编译器拥有更大的编译时间预算, 因而更愿意投入 live range splitting, eviction, coalescing 和更复杂的 cost model.
+
+不过 "Linear Scan 适合 JIT, graph coloring 适合 AOT" 只能当作历史上的大体倾向. 高质量 Linear Scan 可以相当复杂, 图着色 allocator 也可以通过工程设计控制编译成本. 最终还是要看具体实现.
+
+## 7 Linear Scan 和 Graph Coloring 到底差在哪里
+
+经过前几部分以后, 两者现在可以做一个更准确的比较.
+
+给定 liveness 信息, graph coloring 倾向于把它投影成 interference relation. 一旦得到图, allocator 很容易看到一个 value 整个生命周期中和谁冲突, coalescing 也可以通过图收缩来表达. 代价是大量位置信息被压缩掉了, spill placement 和 live range splitting 需要额外机制重新引入这些信息.
+
+Linear Scan 保留的是 interval position. 它很自然地知道一个 value 何时开始, 何时结束, 哪些地方存在 holes, 下一个 use 在哪里. 这让局部 splitting 和 position-based decisions 很直观. 它的 allocation 顺序受线性扫描结构约束更强, 已经做出的早期选择可能阻碍后面的高价值 interval. 如果允许随意回头撤销这些 assignment, 算法就逐渐失去最初简单的 scan 结构.
+
+可以用一个例子来看这种差别. 假设扫描早期遇到一个低价值的长 interval $a$, 当时有空闲寄存器, Linear Scan 很自然地把它分配给 $R_0$. 后面遇到一个位于热循环中的高价值 interval $b$, 却发现所有寄存器已经被早期 assignment 占据. allocator 此时需要 spill, split 或者驱逐旧 interval.
+
+Graph-based 或 priority-based allocator可以先根据 spill weight 处理高价值 interval $b$, 再安排 $a$. LLVM 3.0 的新 allocator正是放弃了 Linear Scan 固定的线性访问顺序, 改用 priority queue 按 spill weight 等优先级处理 live ranges, 并使用 per-physical-register live interval unions 做 interference query. 这样 allocator 可以先处理重要的长 live ranges, 并在后面发现更合适的 assignment 时 eviction 已经分配的低权重 ranges.
+
+## 8 应该怎样理解 Linear Scan 的地位
+
+Linear Scan 的价值不应只用 "快但代码质量差" 来概括. 最基础的版本确实非常简单, 但 interval splitting, lifetime holes, inactive sets, fixed intervals, next-use heuristics 和 register preferences 都可以逐步加入. 一套成熟的 Linear Scan allocator 可以生成相当好的代码, 其核心优势是 live interval 的位置结构天然保存在算法里.
+
+Graph coloring 的经典优势则在于它拥有全局 conflict topology. Chaitin-Briggs 可以在整个 interference graph 上判断 degree, simplify 和 coalescing, 不受到单一 start-position scan order 的严格限制. 但显式图本身也带来构造和维护成本, 而且 spill placement 仍然需要回到程序位置上处理.
+
+现代工业 allocator 往往吸收两边的思想. LLVM Greedy 使用 LiveIntervals, 不建立 textbook Chaitin 式完整冲突图; 它又允许 arbitrary allocation order, eviction 和 extensive splitting, 因而脱离了经典 Linear Scan 的单向扫描限制. 后面学习 LLVM Greedy 时, 会看到 `LiveInterval`, `LiveRegMatrix`, priority queue, eviction 和 `SplitEditor` 怎样组合成这种结构.
+
+从目前六部分的内容看, 三条路线已经可以放在同一张图里:
+
+```text
+                         Liveness
+                            |
+             +--------------+--------------+
+             |                             |
+             v                             v
+     Interference Graph              Live Intervals
+             |                             |
+             v                             v
+     Chaitin / Briggs               Linear Scan
+             |                             |
+     global graph view              scan-order view
+             |                             |
+             +--------------+--------------+
+                            |
+                            v
+                Industrial hybrid designs
+                  e.g. LLVM Greedy
+```
+
+下一部分进入 SSA 与寄存器分配. 前面我们暂时把 live ranges 当成一般结构, SSA 会给它们增加很强的 dominance 性质. 这会影响 interference graph 的结构, phi 的 liveness 语义, coalescing, parallel copy 以及 SSA destruction, 也会解释为什么一些在一般 interference graph 上很困难的问题, 到 SSA 形式下会表现出更特殊的性质.
diff --git a/content/notes/register-allocation/liveness/index.md b/content/notes/register-allocation/liveness/index.md
new file mode 100644
index 0000000..e88dba4
--- /dev/null
+++ b/content/notes/register-allocation/liveness/index.md
@@ -0,0 +1,257 @@
+---
+title: "第二部分: 活跃性分析与生命周期"
+slug: "liveness"
+lang: zh-Hans
+series: register-allocation
+weight: 30
+created: 2026-08-17
+updated: 2026-09-09
+license: CC-BY-SA-4.0
+---
+
+寄存器分配面对的是一种随程序位置变化的资源需求. 一个函数里可能出现几百个 virtual registers, 但 allocator 并不需要同时为它们全部准备物理寄存器. 某个 value 在完成最后一次有意义的 use 之后, 保存它的寄存器就可以重新利用. 因此, 在开始讨论 interference graph 或 Linear Scan 之前, 首先需要确定每个 value 在程序的哪些位置仍然必须保持有效. 这就是 liveness analysis.
+
+## 1 一个 value 在什么时候是 live 的
+
+考虑一段简单的直线代码:
+
+```text
+1: v1 = load a
+2: v2 = v1 + 1
+3: v3 = v2 * 2
+4: return v3
+```
+
+第 1 条指令产生 `v1`. 在第 1 条和第 2 条之间, `v1` 必须被保存, 因为第 2 条指令还要读取它. 第 2 条执行完以后, 后续已经没有 `v1` 的 use, 因而 `v1` 可以死亡. 从这个位置开始, 覆盖 `v1` 所使用的物理寄存器不会再改变程序结果.
+
+更严格地说, 对某个程序位置 $p$ 和 value $v$, 如果从 $p$ 出发存在一条可执行的控制流路径, 路径后面会读取当前的 $v$, 并且在这次读取之前没有新的 definition 取代它, 那么 $v$ 在 $p$ 处是 live 的. 这个定义包含 "存在一条路径" 这一条件, 所以 liveness 属于 may analysis. 编译器不能因为某条路径执行概率很低就忽略它, 只要该路径在语义上可能发生, 对应的 value 就必须得到保留.
+
+这也解释了为什么 liveness 是 backward dataflow analysis. 判断某个 value 当前是否 live, 要看它未来是否还有 use. 信息传播方向因此和程序执行方向相反. 如果 CFG 中有 `B1 -> B2 -> B3`, 程序沿箭头向前执行, liveness 则从后继 block 向前驱 block 传播.
+
+对一条机器指令来说, operand 通常可以分成 def 和 use. 例如:
+
+```text
+v3 = ADD v1, v2
+```
+
+`v1` 和 `v2` 是 uses, `v3` 是 def. `v1` 和 `v2` 必须在指令读取它们时仍然 live, `v3` 则从这次 definition 开始产生新的值. 如果这条 ADD 是 `v1` 的最后一次 use, 那么这个 use 可以带有 kill 的含义. Kill 描述 value 生命周期在这里结束, 并不对应一条真实的 "删除寄存器" 指令.
+
+反过来, 一个 definition 也可能产生从来没有被使用过的 value. 例如:
+
+```text
+v1 = ADD v2, v3
+v4 = MUL v2, v3
+```
+
+如果之后没有任何指令读取 `v1`, 那么第一条指令对 `v1` 的 definition 是 dead def. 在较早的优化阶段, 这种代码通常会被 DCE 删除, 但机器级变换仍然可能临时产生 dead definitions, 所以后端的 liveness 表示也需要能够处理它们.
+
+这里最好始终把 liveness 理解成 value 的性质, 不要和源语言变量的作用域混在一起. 例如:
+
+```text
+x = 1
+use x
+x = 2
+use x
+```
+
+两次 assignment 虽然都写着 `x`, 对数据流来说却产生两个不同 values. SSA 会把它们直接写成 `x1` 和 `x2`. 即使 IR 当前没有处于 SSA form, liveness 分析真正关心的也仍然是哪一个 definition 产生的值在后面还会被读取.
+
+## 2 从 Basic Block 到数据流方程
+
+在没有分支的代码中, 从函数尾部向前逐条扫描就能得到 liveness. 有了 CFG 以后, 一个 block 的出口可能通向多个 successors, 这时通常先在 basic block 粒度求出边界信息, 再把结果展开到 block 内部.
+
+对 basic block $B$, 先定义 $\mathrm{USE}(B)$ 和 $\mathrm{DEF}(B)$. $\mathrm{DEF}(B)$ 包含 block 内产生的 definitions. $\mathrm{USE}(B)$ 包含那些在 block 内第一次使用时还没有在本 block 中被定义的 values, 也就是 upward-exposed uses.
+
+例如:
+
+```text
+B:
+    v3 = v1 + v2
+    v4 = v3 * v5
+```
+
+这里 `v1`, `v2`, `v5` 都需要从 block 外部带进来, 所以它们属于 $\mathrm{USE}(B)$. `v3` 虽然在第二条指令里被使用, 但它已经在第一条指令中定义, 因此不属于 $\mathrm{USE}(B)$. 对这个 block 有 $\mathrm{USE}(B)={v1,v2,v5}$, $\mathrm{DEF}(B)={v3,v4}$.
+
+随后定义 $\mathrm{LIVE\_IN}(B)$ 和 $\mathrm{LIVE\_OUT}(B)$. 前者表示进入 $B$ 时必须仍然有效的 values, 后者表示离开 $B$ 时仍然必须有效的 values. 对普通 CFG, 经典 liveness 方程是:
+
+$$  
+\begin{aligned}  
+\mathrm{LIVE\_IN}(B) &= \mathrm{USE}(B)\cup\left(\mathrm{LIVE\_OUT}(B)-\mathrm{DEF}(B)\right), \  
+\mathrm{LIVE\_OUT}(B) &= \bigcup_{S\in\mathrm{succ}(B)}\mathrm{LIVE\_IN}(S).  
+\end{aligned}  
+$$
+
+第一条方程可以直接按程序语义理解. 进入一个 block 时需要保留的 value 有两种来源. 一种是 block 自己马上会读取, 且在读取之前没有本地 definition 的 value. 另一种是 block 后面的代码还会使用, 同时当前 block 没有重新定义它的 value. 如果一个 value 在 $\mathrm{LIVE\_OUT}(B)$ 中, 但 $B$ 自己会重新定义它, 那么从 $B$ 入口带入的旧值就不需要为了后续 use 继续保留.
+
+第二条方程来自控制流分叉. 如果 $B$ 有两个 successors $S_1$ 和 $S_2$, 某个 value 只要在其中任意一个 successor 入口需要保持, 离开 $B$ 时就必须保存. 假设 $\mathrm{LIVE\_IN}(S_1)={a,b}$, $\mathrm{LIVE\_IN}(S_2)={b,c}$, 那么 $\mathrm{LIVE\_OUT}(B)={a,b,c}$.
+
+用一个 CFG 走一遍会更清楚:
+
+```text
+          B1
+         /  \
+        v    v
+       B2    B3
+        \    /
+         v  v
+          B4
+```
+
+对应代码为:
+
+```text
+B1:
+    v1 = load a
+    v2 = load b
+    if cond goto B2 else B3
+
+B2:
+    v3 = v1 + v2
+    goto B4
+
+B3:
+    v3 = v1 - v2
+    goto B4
+
+B4:
+    v4 = v3 * 2
+    return v4
+```
+
+暂时假设 `cond` 是进入 `B1` 之前已经存在的 value. 对 `B4`, 返回以后没有继续需要保留的 virtual register, 所以 $\mathrm{LIVE\_OUT}(B4)=\varnothing$. `B4` 在定义 `v4` 之前需要 `v3`, 因而 $\mathrm{LIVE\_IN}(B4)={v3}$.
+
+`B2` 和 `B3` 的唯一 successor 都是 `B4`, 所以它们的 live-out 都是 `{v3}`. 两个 blocks 都会在内部重新定义 `v3`, 同时读取 `v1` 和 `v2`, 因而它们的 live-in 都是 `{v1,v2}`. `B1` 的两个 successors 入口都需要 `v1` 和 `v2`, 所以离开 `B1` 时这两个 values 都必须保留. `B1` 本身定义 `v1` 和 `v2`, 因而进入 `B1` 之前只需要 `cond`.
+
+如果 CFG 没有环, 按适当顺序传播这些集合通常很快就能得到结果. Loop 会形成循环依赖. 一个 loop header 的 live-in 可能影响 back edge predecessor 的 live-out, 后者又继续影响 loop header. 因此编译器通常从空集开始反复应用数据流方程, 直到所有集合都不再变化.
+
+```text
+initialize LIVE_IN and LIVE_OUT to empty sets
+
+repeat:
+    for each block B:
+        new_out = union of LIVE_IN of B's successors
+        new_in  = USE[B] union (new_out - DEF[B])
+
+        update LIVE_IN[B] and LIVE_OUT[B]
+
+until no set changes
+```
+
+这个过程会达到 fixed point. 函数中的 relevant values 数量有限, 每个 live set 也只是这个有限集合的子集. 在标准 liveness analysis 中, 信息从空集开始单调传播, 一个 value 加入某个集合以后, 只有有限种状态可供继续扩展, 所以迭代最终会稳定.
+
+工业实现通常使用 worklist 避免反复扫描完全不受影响的 blocks. 对 backward analysis 来说, 如果某个 block 的 `LIVE_IN` 改变, 真正可能因此需要重新计算的是它的 predecessors. 选择合适的遍历顺序也会影响收敛速度, 但不会改变 fixed point 的语义结果.
+
+## 3 Instruction-level Liveness 与 Live Range
+
+Basic block 边界上的 liveness 只是第一步. 寄存器分配通常还需要知道一条具体机器指令前后有哪些 values live. 得到 $\mathrm{LIVE\_OUT}(B)$ 后, 可以从 block 尾部向前逐条扫描. 对普通指令 $I$, 有 $LiveBefore(I)=Use(I)\cup(LiveAfter(I)-Def(I))$.
+
+考虑:
+
+```text
+1: v1 = load a
+2: v2 = load b
+3: v3 = v1 + v2
+4: v4 = v3 * v1
+5: store v4
+```
+
+从末尾开始, 第 5 条指令之前需要 `{v4}`. 穿过第 4 条指令以后, `v4` 的旧值不需要保留, 但这条指令需要 `v3` 和 `v1`, 所以第 4 条之前是 `{v1,v3}`. 再穿过第 3 条, `v3` 被该指令定义, 而 `v1`, `v2` 是 operands, 所以第 3 条之前是 `{v1,v2}`. 继续向前可以得到完整结果:
+
+```text
+                 live before
+
+1: v1 = load a      {}
+                    |
+                    | {v1}
+                    v
+
+2: v2 = load b      {v1}
+                    |
+                    | {v1, v2}
+                    v
+
+3: v3 = v1 + v2     {v1, v2}
+                    |
+                    | {v1, v3}
+                    v
+
+4: v4 = v3 * v1     {v1, v3}
+                    |
+                    | {v4}
+                    v
+
+5: store v4         {v4}
+                    |
+                    | {}
+```
+
+这里可以直接看到物理寄存器复用发生在哪里. `v2` 在第 3 条指令完成以后已经死亡, 所以它占据的寄存器随后可以交给 `v3`. `v3` 在第 4 条指令读取以后也不再需要, 如果目标指令允许 result 覆盖某个已经死亡的 operand, 还可以进一步复用同一个 physical register.
+
+真实 machine instruction 内部有时还需要更细的位置划分. 普通三地址指令可以粗略看成先读取 uses, 再产生 defs, 但 two-address instruction, tied operand 和 early-clobber 会让 def/use 边界更加敏感. 例如:
+
+```text
+R0 = ADD R0, R1
+```
+
+旧 `R0` 是 input, 新 `R0` 是 output. Allocator 必须知道旧值什么时候读完, 新值什么时候可以开始覆盖同一个物理存储. LLVM 的 `SlotIndex` 就是用更细的 machine positions 表示这种关系, 后面讨论 LLVM LiveIntervals 时会再看它的具体设计.
+
+把一个 value 所有 live 的程序位置收集起来, 就得到它的 live range. 在直线代码中可以简单画成:
+
+```text
+instruction:
+1    2    3    4    5    6
+
+v1:
+     |----------------|
+```
+
+CFG 中的 live range 则是一组控制流位置, 不一定适合看成一根连续的横线. 某个 value 可能只在一个分支上需要保持:
+
+```text
+          B1
+         /  \
+        /    \
+       B2    B3
+       |      |
+     live    dead
+        \    /
+         \  /
+          B4
+```
+
+因此 live range 首先是控制流语义上的对象. 只有把机器程序的 positions 线性编号之后, 才会得到便于 Linear Scan 或其他 allocator 使用的 interval representation.
+
+## 4 Live Interval, Segment 与 Lifetime Hole
+
+假设机器程序已经建立线性位置编号, 某个 value 在 `[4,12)` 和 `[20,28)` 两段位置 live. 可以表示成:
+
+```text
+v1:
+[4, 12)        [20, 28)
+```
+
+这两个连续区域通常称为 segments. 中间不 live 的部分是 lifetime hole. 一个现代 `LiveInterval` 因而可以包含多个 segments, "interval" 这个名字并不意味着它只能表示一个单独的数学区间.
+
+Live range 和 live interval 可以从抽象层次上区分. Live range 表示 value 在 CFG 上所有需要保持的程序位置. Live interval 是 allocator 为这些位置建立的线性化数据结构. 如果算法只保留最早 definition 和最后 use, 把整个生命周期粗略表示成一个 `[start,end)`, 实现会很简单, 但也会丢掉 holes.
+
+例如:
+
+```text
+v1: |--------|          |--------|
+
+v2:           |--------|
+```
+
+`v2` 正好位于 `v1` 的 lifetime hole 中. 两个 values 在这些位置上可以共享物理寄存器. 如果把 `v1` 粗略表示成:
+
+```text
+v1: |----------------------------|
+```
+
+allocator 就会认为 `v1` 和 `v2` 发生了整段 overlap, 从而产生不必要的资源冲突.
+
+Linear Scan 的高级版本因此会维护 multiple segments 和 inactive intervals. LLVM Greedy 同样依赖精细的 live interval information 来判断物理寄存器冲突和选择 splitting points. 与之相对, 经典 Chaitin-style allocator 会进一步把这些位置信息压缩成 interference graph, 只保留两个 allocation objects 是否曾经发生过冲突.
+
+Liveness 也直接决定 register pressure. 如果某个位置同时 live 的 values 是 `{a,b,c,d}`, 并且四个 values 都竞争同一个 register class, 那么这一位置至少产生 4 个寄存器单位的需求. 如果 target 只有三个可用寄存器, allocator 就需要通过 spilling, splitting, rematerialization 或其他变换降低实际寄存器占用. Instruction scheduling 改变 definitions 和 uses 之间的距离时, live ranges 也会随之伸长或缩短, 因而 register pressure 和 RA 从来没有完全脱离代码调度.
+
+下一步把 liveness 转换成 interference relation时, allocator关心的就是哪些 values 曾经在同一个程序位置同时需要相同物理资源. Graph coloring 会把这种关系记录成边, Linear Scan 和 LLVM Greedy 则更直接地利用 live interval overlap. 两种路线的数据结构不同, 但它们都建立在这一部分得到的生命周期信息上.
diff --git a/content/notes/register-allocation/llvm-gcc/index.md b/content/notes/register-allocation/llvm-gcc/index.md
new file mode 100644
index 0000000..1659de6
--- /dev/null
+++ b/content/notes/register-allocation/llvm-gcc/index.md
@@ -0,0 +1,360 @@
+---
+title: "第九部分: CPU 工业寄存器分配: LLVM 与 GCC"
+slug: "llvm-gcc"
+lang: zh-Hans
+series: register-allocation
+weight: 100
+created: 2026-08-18
+updated: 2026-09-09
+license: CC-BY-SA-4.0
+---
+
+前八部分建立的各种概念在 LLVM 和 GCC 中都有对应物, 但它们寄存器分配策略大相径庭. LLVM 的优化型 allocator 长期围绕 `LiveInterval` 工作, 保留 value 在机器代码位置上的生命周期信息, 然后针对某个 physical register 查询实际 interference. 当 assignment 失败时, Greedy 可以 eviction, splitting, recoloring, 最后才进入 spill. GCC 则把主要工作分成 IRA 和 LRA 两层. IRA 构造 allocation regions 和 conflict graph, 使用经过大量工程扩展的 Chaitin-Briggs coloring 做全局分配; LRA 接着根据 RTL instruction constraints 处理 reload, late assignment, splitting 和 rematerialization.
+
+这两种设计恰好代表了前面几部分中的两条不同路线. LLVM Greedy 延续了 live interval 的位置化表示, 但打破了经典 Linear Scan 的单向 allocation order. GCC IRA 则保留了 graph coloring 的基本框架, 同时加入 region hierarchy, irregular register classes, cost propagation 和 live range splitting. 看清两者的内部结构之后, "工业寄存器分配为什么远比 textbook coloring 复杂" 这个问题也会具体很多.
+
+## 1 LLVM 的 Machine Register Allocation 基础设施
+
+LLVM 进入 Machine IR 后, operand 已经逐渐具有目标机器语义. Virtual register 属于某个 `TargetRegisterClass`, instruction operands 可以带有 tied, early-clobber, implicit use/def 等属性, function call 则通过 register mask 描述哪些 physical registers 会被破坏. 在 register allocation 以前, virtual registers 仍然可以利用 Machine SSA 提供的 single-definition 结构; RA 附近的 passes 会逐渐把 phi, two-address constraints, copies 和 physical-register requirements 转换成 allocator 能直接处理的形式.
+
+LLVM 对 virtual register 生命周期的核心表示是 `LiveInterval`. 它建立在 `LiveRange` 上, 一个 `LiveRange` 由按照 `SlotIndex` 排序的多个 `Segment` 组成. Segment 描述一个 value 在 `[start,end)` 上 live, `VNInfo` 则表示其中某个 definition 产生的 value number. 当一个 live range 中出现多个 definitions, CFG join 或 splitting 产生的新 fragments 时, allocator 仍然可以用不同的 `VNInfo` 区分这些值. `LiveInterval` 还可以保存 subregister subranges, 因而同一个 virtual register 内部不同 lanes 的 liveness 也能够进一步细分.
+
+`SlotIndex` 把 Machine IR 中的程序位置映射到一个有序空间. 它的粒度比 "第几条指令" 更细, 可以区分 instruction boundary, register use/def slot 以及 early-clobber def. 因此一条 live interval 不只是抽象地知道自己和另一条 interval 有冲突, 它还保留冲突发生在函数的什么位置. 这份位置信息后来直接用于 local splitting, global splitting, spill placement 和 fixed-register interference.
+
+Physical register 一侧由 `LiveRegMatrix` 组织. 它沿 `SlotIndex` 和 register unit 两个维度记录已经建立的 assignment. Register unit 是处理 overlapping physical registers 时使用的最小 interference unit. 因此 x86 上两个不同的 architectural register names 即使发生 alias, allocator 仍然能够在共享的 register units 上检测到冲突. 这就是第八部分讨论的 subregister aliasing 在 LLVM RA 中落到数据结构后的样子.
+
+假设 virtual register `%v1` 的 live interval 为:
+
+```text
+%v1:
+      |-----------|         |------|
+```
+
+allocator 准备尝试 physical register `R0`. `LiveRegMatrix` 会查看 `R0` 及其 aliases 覆盖的 register units, 再检查 `%v1` 的 segments 是否与这些 units 上已经存在的 live ranges 重叠. 如果没有重叠, `%v1 -> R0` 可以直接建立. 如果冲突来自另一个已经分配的 virtual register, 那个 assignment 仍然可能被 eviction 或 recoloring 改变. 如果冲突来自 fixed register-unit live range, 或 `%v1` 跨越了一个不 preserve `R0` 的 register mask, 问题就不能通过简单驱逐另一个 virtual register解决. `LiveRegMatrix` 因此区分 free, virtual-register interference, fixed register-unit interference 和 regmask interference.
+
+这一结构和经典 interference graph 有明显差别. Chaitin-Briggs 会先建立 virtual nodes 之间的冲突边, 随后主要在图结构上工作. LLVM Greedy 给 `%v1` 尝试 `R0` 时, 直接询问 "`R0` 对 `%v1` 的这些 live positions 是否可用". Interference 仍然存在, 只是它没有被统一压缩成一张完整的 virtual-register graph. 这让 allocator 在发现某个 register 只有局部位置发生冲突时, 很自然地继续考虑 splitting.
+
+LLVM 的优化型 RA pipeline 也围绕这些数据结构组织. 省略 target-specific 插入点和与本文无关的 passes 后, 主干可以用如下示意图表示.
+
+```text
+Machine SSA
+    |
+    v
+PHI Elimination
+    |
+    v
+Two-Address Instruction Lowering
+    |
+    v
+Register Coalescer
+    |
+    v
+Independent Subregister Renaming
+    |
+    v
+Pre-RA Machine Scheduler
+    |
+    v
+Register Allocator
+    |
+    v
+VirtReg Rewriter
+    |
+    v
+Stack-slot / copy / post-RA optimizations
+```
+
+`PHIElimination` 把 phi 的 edge semantics 转换成后续 machine representation, `TwoAddressInstruction` 处理 two-address 和 tied operand 产生的问题. `RegisterCoalescer` 在正式 allocation 前先合并能够安全合并的 COPY-related live ranges. `MachineScheduler` 随后改变机器指令顺序, 它的选择会直接改变 live range shape 和 register pressure. RA 完成后, `VirtRegRewriter` 根据 `VirtRegMap` 把 MachineInstr 中剩余的 virtual-register operands 改写成 physical registers.
+
+LLVM 同时提供 Fast, Basic, Greedy 和 PBQP 几种 allocator. Fast 面向低优化编译, 重点是降低 allocation 开销; Greedy 是优化编译的默认 allocator. Basic 提供较简单的全局 allocation 框架, Greedy 在相同基础设施上加入更完整的 eviction 和 global live range splitting. PBQP 则把 register allocation 建模成 Partitioned Boolean Quadratic Programming 问题, 通过 PBQP solver 求解 assignment. 个人很欣赏 PBQP, 因为它的数学理论足够优美, 读者可以参考原始论文, 不过后面的讨论集中在 Greedy, 因为它最能体现 LLVM CPU register allocation 的工业设计.
+
+## 2 LLVM Greedy: 从 Priority Queue 到 Eviction
+
+LLVM Greedy 不按照 live interval 的起始位置从前向后扫描. 待分配 intervals 存放在 priority queue 中, allocator 每次从队列中取出一个 live interval进行 assignment. 这里的 priority 也不能简单等同于 spill weight. 对原始的 single-block local ranges, 默认策略保留接近 instruction order 的顺序; 对 global ranges 和 split 后产生的 ranges, 主要按照 live-range size 从长到短处理. Priority 中还可以编码 register class 的 allocation priority, global/local 属性和 physical-register preference.
+
+这种 allocation order 解决了经典 Linear Scan 的一个结构限制. 假设有一条很长的 global interval `a`, 以及大量短 local intervals:
+
+```text
+a:
+|------------------------------------------------------|
+
+b:      |---|
+c:            |---|
+d:                  |---|
+e:                        |---|
+f:                              |---|
+```
+
+Linear Scan 按 start position 前进时, 前面已经完成的 assignment 会逐渐形成既定状态. Greedy 可以优先处理 global 和较长 ranges, 让影响范围大的 assignment 先占据合适的 physical register, 再让局部 ranges 在剩余空间中安排. 对单 block 的简单 SSA ranges, instruction order 本身又具有很好的局部性质, 所以 LLVM 没有机械地用同一个优先级公式处理所有 intervals.
+
+取出一个 interval 后, allocator 先根据 register class, target register order 和 register hints 建立 `AllocationOrder`. Register hint 通常来自 COPY relation 或其他 machine preference. 如果 `%v1` 与一个已经位于 `R3` 的 value 存在有利于 coalescing 的关系, `R3` 就可能被提前尝试. 对每个 candidate physical register, allocator通过 `LiveRegMatrix` 检查 interference. 找到 free candidate 后就可以在 `VirtRegMap` 和 matrix 中建立 assignment.
+
+Physical registers 也有 cost. 第八部分讨论过 callee-saved register 第一次在函数中使用时可能引入 prologue/epilogue save/restore. Greedy 在选择 candidate 和 eviction strategy 时会考虑这种代价, 因此 "某个寄存器目前空闲" 并不自动意味着立即使用它最划算. 当一个便宜 register 与已有 virtual intervals 冲突, 驱逐这些 intervals有时反而比开始使用一个新的高成本 register 更合适.
+
+Eviction 是理解 LLVM Greedy 的核心机制之一. 假设 `R0` 已经分配给 `b`, 新到来的 `a` 与它发生 overlap:
+
+```text
+R0:
+
+b:          |----------------|
+a:       |------------------------|
+             <---- conflict ---->
+```
+
+如果 allocator认为 `a` 更值得占据 `R0`, 可以撤销 `b -> R0`, 再建立 `a -> R0`. `b` 不会因此立即 spill, 它被重新送回 allocation 流程, 后面仍然可以获得别的 physical register, 被 split, 或最终 spill. 一次 eviction 还可以同时驱逐同一个 candidate register 上的多条 interfering intervals. Greedy 会根据 victim costs, spillability 和 cascade information 控制这种行为, 避免 allocations 在几个 intervals 之间无限来回震荡.
+
+这和 Chaitin-Briggs 中的 stack coloring 形成了不同的搜索方式. Chaitin-Briggs 先通过 simplify 得到一个 coloring order, 再在 select 阶段恢复节点. Greedy 在 assignment 过程中允许局部撤销已有结果. 当一个高价值 live range 到来时, allocator 可以重新安排先前的较弱 assignment, 因此 allocation state 本身处于持续变化中.
+
+Spill weight 在这里主要作为成本信息参与 eviction, splitting 和 spill decision. `LiveInterval` 的 weight 会根据 uses, defs 和 block frequency 等信息计算, 反映保留该 interval 的价值. 但 queue priority 和 spill weight 是两个不同概念. 一个 global interval 可以因为长度和 allocation stage 很早进入队列, 它是否值得驱逐其他 intervals 或最终进入 memory, 则要结合另一套 cost information 判断.
+
+## 3 LLVM Greedy: Splitting, Recoloring 与 Spilling
+
+如果一条 live interval 找不到 free physical register, eviction 也没有得到合适解, Greedy 不会立即把整个 interval spill. 第一次 allocation failure 会把它转入 split stage 并重新排队. 其他更容易处理的 intervals 继续完成 assignment, 等这条 interval 再次进入 allocator时, 周围的 physical-register occupancy 已经稳定了许多. 此时 allocator可以更准确地看到 "冲突究竟集中在哪些程序区域", 再决定怎样 splitting.
+
+考虑一条 global live interval:
+
+```text
+v:
+|-----------------------------------------------------------|
+
+R4 interference:
+       |------|                 |---------|
+```
+
+如果 `R4` 在绝大部分生命周期中都很适合 `v`, 只有两个局部区域发生 interference, 把整个 `v` 放弃掉会损失很多 register residency. Greedy 可以把原始 interval 切成多个 child intervals, 让其中适合 `R4` 的部分继续使用 `R4`, 剩余 fragments 回到队列重新寻找 assignment:
+
+```text
+v0: |------|
+
+v1:        |--------------------|
+
+v2:                             |---------|
+
+v3:                                       |---------------|
+```
+
+Splitting 不只是修改一个 interval 的 `[start,end)` 元数据. `SplitKit` 和 `SplitEditor` 会创建新的 virtual registers 和 live intervals, 重新建立 value numbers, 修改 fragment 与 Machine IR uses 的对应关系, 并处理进入或离开某个 child interval 时需要的数据传递. Split mode 还可以偏向减少动态 COPY cost 或减少静态 COPY 数量, 因此 split boundary 的选择本身就是 code-quality decision.
+
+Global splitting 会结合 basic-block frequency 和 use blocks. 假设一个 value 在热循环中频繁使用, 在冷控制流中偶尔存活, allocator可能让热区域保持 register assignment, 把 transfer 或 spill pressure 推向较冷路径:
+
+```text
+cold          hot loop                cold
+
+------|=========================|---------
+      ^                         ^
+   transfer                  transfer
+```
+
+这种策略体现了 position-based allocator 的优势. Interference 发生在 live interval 的哪些 blocks, 哪些 edge 上需要进入或离开 register region, 都直接存在于 allocator使用的数据结构里. Greedy 可以围绕一个 candidate physical register 的具体 interference pattern选择 split region.
+
+Splitting 后生成的 child intervals重新进入 queue, 所以后续过程依然是普通 register allocation. 某个 child 可能立即找到 free register, 可能 eviction 其他 intervals, 也可能继续 split. 这使 "一个 virtual register 的 allocation" 不再是一次性决策. 原始 live range 可以逐渐被分解成一组具有不同 assignment 的 fragments.
+
+Greedy 还保留 recoloring 作为晚期搜索手段. 当某个 interval无法获得物理寄存器时, allocator可以尝试改变周围 interfering virtual registers 的 assignment, 递归寻找另一组合法组合. Recoloring 只能够重新安排 virtual-register interference; fixed register-unit interference 和 regmask clobber 属于机器已经给定的障碍, 不能通过重新着色消失. 为了限制编译时间, recoloring 的深度和搜索规模会受到控制.
+
+最终仍然无法安排的 interval 进入 spiller. `InlineSpiller` 会插入 spills 和 reloads, 同时尝试 stack-access folding 和 rematerialization. 如果某个 definition 足够便宜, use 位置可以重新生成这个值, allocator就没有必要把它先 store 到 stack 再 load 回来. Spill rewrite 还可能产生新的短 virtual registers, 这些 fragments继续接受后端后续处理.
+
+把 Greedy 的主要控制流放在一起, 可以看到它与 textbook allocator 的结构差异:
+
+```text
+priority queue
+      |
+      v
+try free physical register
+      |
+      +---- success ----------------------> assign
+      |
+      v
+try eviction
+      |
+      +---- success ----------------------> assign
+      |                                      |
+      |                                      +--> requeue victims
+      v
+defer and requeue
+      |
+      v
+try splitting
+      |
+      +---- create child intervals -------> requeue children
+      |
+      v
+try recoloring
+      |
+      +---- success ----------------------> assign
+      |
+      v
+spill / rematerialize
+```
+
+Register coalescing 与 Greedy 的关系也因此比 IRC 更分散. LLVM 有独立的 `RegisterCoalescer` 在 RA 前合并 live ranges, Greedy 自身还会使用 register hints, hint-related splitting 和 recoloring 尽量保持有利的 COPY assignment. Copy elimination 并没有集中在一个 `Simplify-Coalesce-Freeze` 状态机中, 它贯穿 coalescing pass, allocation order 和后续 repair decisions.
+
+## 4 GCC IRA: 从 RTL Pseudo 到 Regional Graph Coloring
+
+GCC 在机器级主要使用 RTL. Register allocation 以前, RTL 中大量值仍然由 pseudo-registers 表示, 最终需要落到 hard registers 或 memory. IRA 的 allocation entity 称为 `allocno`. 一个 allocno 表示某个 pseudo-register 在一个 allocation region 中的生命周期, 因而同一个 pseudo 可以在不同 nested regions 中对应不同 allocnos.
+
+IRA 的 regions 形成一棵树. Root region 覆盖整个函数, 其他主要 regions 来自 natural loops. 如果函数中存在两层嵌套循环, 结构可以近似表示为:
+
+```text
+function region
+|
++-- allocno(p)
+|
++-- loop L1
+|   |
+|   +-- allocno(p)
+|   |
+|   +-- loop L2
+|       |
+|       +-- allocno(p)
+|
++-- other code
+```
+
+外层 allocno 会累积下层 regions 中与同一 pseudo 相关的 cost, conflict, copy 和 hard-register information. 因而 root region 的 coloring 能够先形成整个函数范围的 global decision, 随后进入 loops 和 subloops 时再根据局部情况改善 assignment. 这种 top-down regional allocation 同时保留 global view 和 loop-local optimization 空间.
+
+Allocno 内部还可以分解成更细的 `ira_object`. `ira_object` 保存 conflict information 和 live ranges, 并且可以对应 allocno 的某个 subword. 对一个 multi-word allocation object, IRA 因而能够在比整个 pseudo 更细的层面描述 conflicts. Live range 使用整数 program points 表示, 这些 points 位于 operand die 和 output born 等可能改变 liveness 的机器位置附近. 两个 objects 的 live ranges 相交时, IRA 据此建立 conflict relation.
+
+IRA 在 coloring 前会同时构造几类信息. Allocno class 决定某个 allocno 可以使用哪一类 hard registers, pressure class 用于计算不同硬件资源上的 register pressure. 每个 allocno 还拥有 memory cost, hard-register cost vector, conflict hard-register costs, crossed-call 信息以及与其他 allocnos 的 copy relation. 跨 call 的 allocno 会让某些 hard-register choices 增加 save/restore cost; move-related allocnos 则可以通过修改 hard-register preference, 增加获得相同 register 的机会.
+
+因此 IRA 的输入已经远远超过一张无权 interference graph. 可以用如下示意图简要概括.
+
+```text
+RTL pseudos
+    |
+    v
+region tree
+    |
+    v
+allocnos / objects
+    |
+    +--> live ranges
+    +--> allocno conflicts
+    +--> hard-register conflicts
+    +--> register pressure
+    +--> memory costs
+    +--> hard-register costs
+    +--> copy preferences
+    |
+    v
+regional graph coloring
+```
+
+IRA 名字中的 "Integrated" 也可以从这里理解. Coalescing, hard-register preference 和 live range splitting 都被放进 regional coloring 的整体过程. Copy relation 会改变 coloring 时的 register preference, regional allocation 的不同结果又可以自然产生 live-range splitting.
+
+IRA 的 coloring 仍然能够清楚看到 Chaitin-Briggs 的血统. Allocnos 被逐步压入 coloring stack, high-pressure 情况使用 Briggs optimistic coloring, 不会在入栈阶段就把所有 high-degree nodes立即判定为 actual spill. 弹栈时再根据已经占据的 hard registers 尝试 assignment. 如果一个 coalesced allocno无法找到合法 hard register, IRA 还可以撤销 coalescing, 将拆开的 allocnos重新送回 coloring process.
+
+IRA 还会形成 `thread`. 一个 thread 由互不冲突且通过 copies 联系起来的 colorable allocnos 组成. 把这些 allocnos相邻压入 coloring stack, 后续分配同一个 hard register 的机会更高, 从而减少 move. 这和第五部分 IRC 追求的目标相同, 但 IRA 没有采用 `Simplify`, `Coalesce`, `Freeze` 那套 textbook worklists, copy preference 被嵌入 regional coloring 和 hard-register cost 中.
+
+第八部分讨论的 irregular register classes 会直接修改 IRA 的 trivial-colorability 判断. 假设某个 allocno可以使用多个 general hard registers, 它同时和 8 个 allocnos发生 conflict, 但那 8 个 allocnos都只能使用 `EAX`. 只看 graph degree 会得到 8 个 neighbors, 可这些 neighbors 最多共同封锁 `EAX` 这一种 hard-register choice, 其余 general registers仍然可以使用. IRA 会结合 allocno classes 和可分配 hard-register sets 判断一组 conflicts 实际能够排除多少候选资源.
+
+因此 GCC 虽然使用 coloring stack 和 optimistic coloring, 但颜色集合已经具有 register-class structure, 每种 hard register 还有不同 allocation cost, copies 修改 preference, regions 又让同一个 pseudo 在不同程序区域拥有不同 allocation objects. GCC 公开的 `-fira-algorithm=CB` 对应 Chaitin-Briggs coloring, 另一个选择是 priority coloring; speed-oriented regional allocation通常使用 loop regions, 并可以过滤 register pressure 较低的 loops.
+
+Regional allocation 本身还会产生 splitting. 假设同一个 pseudo `p` 在外层代码和热循环中分别对应两个 allocnos, IRA 可能得到:
+
+```text
+outside loop          hot loop           outside loop
+
+     R2      --->        R5       --->       R2
+```
+
+如果 loop 内使用 `R5` 的综合成本更低, IRA 可以在 region boundary 创建新的 pseudo 并插入 transfer code. 这样 `p` 的逻辑生命周期就被拆成多个具有不同 location 的 fragments. Splitting 在这里来自 region hierarchy, 与 LLVM 根据 physical-register interference pattern 主动寻找 split region 的方式不同.
+
+IRA coloring完成后还会从 cost 角度继续改善结果. 某些 allocnos即使可以获得 hard register, memory assignment 也可能使整体成本更低; 某些已经 spill 的 allocnos释放出的 registers又可能让其他高价值 allocnos获得更便宜的 assignment. 因此 "尽可能多地 color nodes" 并不等价于 "得到最低-cost machine code". IRA 的 hard-register costs, memory costs 和 region-border move costs共同参与这些取舍.
+
+## 5 GCC LRA: Instruction Constraints 与 Reload
+
+IRA 做出的是全局 allocation decision, 真实 RTL instruction 还必须逐条满足 machine description 中的 operand constraints. 某个 pseudo 从全局角度适合放进 GPR, 某次具体 use 却可能要求一个更窄 register class; 某条 instruction alternative 可能允许 memory operand, 另一条要求 hard register; tied operand, early-clobber 和 address constraint 还会进一步限制局部 assignment. LRA 接在 IRA 后面处理这些问题. 它继承了过去 Reload pass 所承担的角色, 但组织方式是围绕 instruction constraints 反复修正 RTL.
+
+假设 IRA 后有:
+
+```text
+OP p
+```
+
+而 `OP` 的合法 alternative 要求这个 operand 位于某个 restricted hard-register class. 如果 `p` 的 location 无法直接满足约束, LRA 可以生成一个 reload pseudo:
+
+```text
+r <- p
+OP r
+```
+
+`r` 拥有符合 instruction constraint 的 allocno class, 后续 hard-register assignment 再给它寻找合法 register. 如果 `p` 已经位于 stack, 这就是经典意义上的 reload temporary. 如果 instruction 支持合适的 memory alternative, LRA 也可能选择另一种 encoding, 从而避免显式 reload.
+
+LRA 的 constraint pass 需要在一条 instruction 的多个 alternatives 中做选择. 一个 alternative 可能减少 reload, 但消耗更稀缺的 register class; 另一个允许 memory, 却可能产生更昂贵的 machine instruction. 选择 alternative 后如果生成新的 reload pseudos, 这些 pseudos又需要 hard registers. 为它们腾出位置可能迫使其他 pseudos spill, 新的 spill location 又可能改变地址表达式. Address displacement 发生变化以后, 原先合法的 address constraint 还可能失效. 因而 LRA 的求解天然带有反馈, 必须迭代直到 instruction 和 address constraints 都得到满足.
+
+可以把这种反馈关系画成:
+
+```text
+instruction constraints
+        |
+        v
+choose alternatives
+        |
+        v
+create reload insns / reload pseudos
+        |
+        v
+hard-register assignment
+        |
+        v
+spill / split / inheritance
+        |
+        v
+stack locations and addresses
+        |
+        +----------------------+
+                               |
+                     constraints changed
+                               |
+                               +----> iterate
+```
+
+第一次 constraint processing 会遍历 instructions 并选择 alternatives. 后续迭代会尽量保留仍然合法的选择, 把工作集中在受到新 reload, spill 或 address change 影响的部分. 这和重新进行一次完整 IRA coloring 的成本结构不同, LRA 的目标是逐步把已经相当接近机器代码的 RTL 修正到可编码状态.
+
+LRA 还有 inheritance optimization. 假设 spilled pseudo `p` 第一次使用时已经 reload 到某个 hard register:
+
+```text
+reload1 <- p
+use reload1
+
+...
+
+reload2 <- p
+use reload2
+```
+
+如果两次 use 之间 `p` 的值没有改变, 第一条 reload 得到的值可能仍然保存在 register 中. LRA 可以创建 inheritance pseudo, 让后一次 use 尝试继承已经存在的 register value, 从而避免再次访问 stack. 后面的 assignment 如果无法给 inheritance pseudo安排合适的 hard register, 这次 transformation 可以撤销, 恢复原来的 reload sequence.
+
+Inheritance pass 同时会进行 EBB 范围内的 live-range splitting. 一个 global pseudo 穿过高 pressure 区域时, LRA 可以在局部插入 save/restore-style transfer, 把原来的长 live range拆开. 跨 call 的 pseudos也可以利用类似方式缩短需要保持特定 hard register 的区域. 这意味着 GCC 的 splitting 不只存在于 IRA region boundaries. IRA 处理全局和 loop hierarchy 上的分割, LRA 还会在更靠近 instruction constraints 的阶段继续调整生命周期.
+
+Rematerialization 则处理另一类 spill cost. 对一个 spilled pseudo, 如果它的 definition可以用较便宜的 machine operation重新计算, LRA 可以在 use 附近重新生成 value, 省掉 memory reload. 候选通常需要避开 memory access 和昂贵的 div/mod 等操作, 同时保证输入 operands 已经处于可用 hard registers 中. GCC 的 `-flra-remat` 就控制这类 CFG-sensitive rematerialization.
+
+因此 LRA 虽然名称中有 "Local", 它处理的范围远不止 basic-block-local register assignment. 它拥有 iterative constraint solving, reload-pseudo creation, hard-register assignment, EBB inheritance, splitting 和 CFG-sensitive rematerialization. 它在 GCC RA 架构中的位置可以理解成从 "IRA 已经给出一个良好的全局资源方案" 走到 "每一条 RTL instruction 都满足真实 ISA constraints" 的机器合法化阶段.
+
+## 6 LLVM Greedy 与 GCC IRA+LRA 的结构差异
+
+把两套 allocator 放在一起以后, 最值得比较是它们怎样保存信息以及怎样从失败的 assignment 中恢复. 如下是一个简要的表格对比二者在重要问题上的差异.
+
+|问题|LLVM Greedy|GCC IRA + LRA|
+|---|---|---|
+|主要 allocation object|`LiveInterval` 及 split fragments|IRA `allocno`, conflict 可以细化到 `ira_object`|
+|生命周期表示|`SlotIndex` 上的 `Segment` 和 `VNInfo`|IRA program points 上的 live ranges|
+|Conflict 组织|`LiveRegMatrix` 按 physical register 和 regunit 查询|IRA 构造 allocno/object conflicts|
+|全局分配方式|priority-based incremental assignment|regional Chaitin-Briggs coloring|
+|已有 assignment 的调整|eviction 和 recoloring|optimistic coloring, coalescing undo, regional reassignment|
+|Coalescing|独立 RegisterCoalescer + hints + recoloring|copies 和 hard-register preference 集成到 IRA coloring|
+|Splitting|根据 interval interference 做 local/global splitting|IRA region splitting + LRA EBB splitting|
+|Machine constraints|MachineInstr constraints, register classes, regunits, regmasks贯穿 RA|IRA 先做全局 allocation, LRA 再集中满足 RTL alternatives 和 address constraints|
+|Spill 后处理|InlineSpiller, folding, rematerialization|LRA reload pseudos, inheritance, splitting, rematerialization|
+
+LLVM 的主要优势在于 allocator一直保留精细的程序位置信息. 当 `%v` 无法使用 `R5` 时, Greedy 可以看到 `R5` 究竟在哪些 segments 上被什么对象占据, 然后决定 eviction 整条 interfering interval, 围绕局部 interference splitting, 或把某个 child fragment 留在 memory. `LiveRegMatrix` 又让同一套查询自然处理 register aliasing 和 call regmask.
+
+GCC IRA 的主要全局视角来自 conflict graph 和 region hierarchy. Allocno 把 "一个 pseudo 在一个 region 中的生命周期" 变成 coloring entity, nested loops使 allocator能够在整个函数方案之上继续改善 hot regions. Chaitin-Briggs 的 simplify/select 思想仍然存在, 但 trivial colorability 已经扩展到 intersected register classes, coloring choice 又受到 hard-register cost vectors 和 copy preferences控制.
+
+IRA 与 LRA 的分工还体现出 GCC 对 machine constraints 的处理层次. IRA 可以在较稳定的全局模型上处理 pressure, conflicts 和 region costs, 避免让每条 instruction alternative 都进入 graph-coloring problem. LRA 在全局方案之后处理高度 target-specific 的约束, 并通过 reload pseudos 和迭代 repair 收敛到合法 machine code. LLVM 则让 target register classes, physical-register interference, regmask 和 MachineInstr operand constraints更早地进入统一的 Machine IR allocation infrastructure.
+
+Scheduler 与 register pressure 在两边也有直接联系. LLVM 的 optimized RA pipeline 在 Greedy 前执行 machine scheduling, 调度后的 live intervals直接成为 allocation 输入. GCC 也支持 pressure-sensitive pre-RA scheduling, `-fsched-pressure` 会在调度时控制 register pressure, 避免调度把同时 live 的值数量推过可用 hard-register capacity, 从而给后续 IRA/LRA 制造额外 spills.
+
+从前几部分的算法谱系来看, LLVM Greedy 很难简单放进 "Linear Scan" 或 "Graph Coloring" 其中一格. 它使用 `LiveInterval` 和 position-based interference, 延续了 interval allocator 的基础设施; allocation order 已经变成 priority queue, 同时加入 eviction, global splitting 和 recoloring. GCC IRA 与 Chaitin-Briggs 的继承关系更加直接, 但 nested regions, irregular register classes, copies, hard-register costs 和 LRA 又把 textbook coloring 扩展成了一个完整的机器级 allocation system.
+
+这两套实现最终都在解决同一组资源问题. 一个 value 应该在哪些区域占据 register, 哪个 physical register 对它最便宜, 哪些 COPY 值得消除, 哪些 conflicts 值得通过重新安排其他 values 解决, 哪些区域应该 split, 什么时候 memory traffic 比继续争夺寄存器更划算. LLVM 把这些决策更多组织在 live interval 和 incremental assignment 周围, GCC 把它们更多组织在 regional coloring 和 constraint repair 周围. 下一部分进入 GPU 后, 这些问题还会多出 occupancy 这一层反馈: 一条 live range 多占用几个 registers, 可能改变整个 SM/CU 同时能够驻留的 warps 或 waves 数量, register allocation 的目标函数也会随之发生变化.
diff --git a/content/notes/register-allocation/preface/index.md b/content/notes/register-allocation/preface/index.md
new file mode 100644
index 0000000..8844d29
--- /dev/null
+++ b/content/notes/register-allocation/preface/index.md
@@ -0,0 +1,93 @@
+---
+title: "前言"
+slug: "preface"
+lang: zh-Hans
+series: register-allocation
+weight: 10
+created: 2026-08-09
+updated: 2026-09-09
+license: CC-BY-SA-4.0
+---
+
+最近一直在进行的工作就是推进我的新编译器, 虽然在现在这个时代一个人实现一个功能完备性能又很好的大型软件几乎不可能, 但是我目前至少在让后端尽可能支持更多的功能以及对代码有更好的优化. 完整的代码暂时没有公开, 但是欢迎各位去看我的生成代码质量比较好的竞赛编译器: [rcc-project](https://github.com/Richard-Qin-X/rcc-project)
+
+说到后端, 寄存器分配无疑是最重要的环节之一, 但是翻了一下知乎, 没有发现系统地讲解寄存器分配的文章, 而且我对编译原理这门课也有一些意见, 花大量篇幅讲代码怎么到 AST 在现在的意义并不大, 中后端的内容太少而且过于 textbook, 与现实脱节太严重, 尤其是深度学习的兴起, MLIR 和 异构平台代码生成更重要了. 于是决定写一系列文章, 系统性地介绍一下**寄存器分配的基本概念, 经典算法, 现代机器模型, CPU 工业实现, GPU 寄存器分配**, 以及一个 **综合与实践** 环节 (一个自己动手实现简单寄存器分配算法的教程).
+
+具体的结构等这个系列写完之后我会放在这篇文章里. 
+
+相信通过这个系列的文章, 读者能够比较好地理解活跃性分析与生命周期, 冲突图, 图着色寄存器分配, 贪心分配, 以及现代编译器 LLVM 和 GCC 的后端架构和重要算法, 还有 GPU 的一些知识.
+
+我觉得最值得期待的就是 **综合与实践** 部分, 毕竟光看理论和自己动手还是不一样的. 我已经在我的新编译器里手写了一个 PBQP, 当然在这个部分我不会介绍这种几千行代码的复杂算法, 我会尽量设计一个比较简单的场景和算法, 让读者能用几百行代码完成.
+
+还有一些比较值得关注的就是 **GPU 寄存器分配**, 传统编译原理课程中讲的寄存器分配主要是针对 CPU, 而 GPU 寄存器分配与 CPU 在约束和优化目标上有一些明显不同. 很多人会疑惑的一点就是 GPU 寄存器这么多分配就不重要了吧, 然而事实上只要存储速度跟不上计算速度寄存器分配就永远是焦点, GPU 上尤其特殊的一点是, 寄存器分配直接影响 occupancy.
+
+假设一个 SM 有固定数量的寄存器
+
+```
+SM register file
+      │
+      ├── Block 0
+      │    ├── Thread 0: 64 regs
+      │    ├── Thread 1: 64 regs
+      │    └── ...
+      │
+      ├── Block 1
+      └── ...
+```
+
+如果每个线程需要的寄存器越多, 一个 SM 同时能够驻留的线程 / warp / block 就可能越少. 比如粗略地说
+
+```
+SM 有 65536 个寄存器
+
+每线程 32 registers
+→ 理论上可容纳 65536 / 32 = 2048 threads
+
+每线程 64 registers
+→ 理论上只可容纳 1024 threads
+```
+
+当然实际 occupancy 还同时受到 warp 数, block 数, shared memory 等硬件限制.
+
+所以 GPU register allocator 往往面对一个很有意思的权衡
+
+```
+多用寄存器
+   ↓
+更少 spill，单线程执行可能更快
+   ↓
+但 occupancy 可能降低
+
+少用寄存器
+   ↓
+occupancy 可能提高
+   ↓
+但可能产生 spill
+   ↓
+访问 local memory，代价可能很高
+```
+
+这里的 local memory 不是片上寄存器旁边的一小块高速内存, 在 CUDA 中，它通常属于显存层次, 虽然可能被 cache, 因此 register spilling 往往比较昂贵.
+
+另外在本篇文章结束前聊一个比较有意思的发现, 现在没有人专门做GCC和LLVM寄存器分配的 benchmark 对比. 而且这件事有点反常, register allocation 学术论文非常多, 但 production allocator vs production allocator 的标准化 benchmark 基础设施却很弱. 我个人觉得一方面的原因是二者后端根本不接受一种形式的 IR, 导致没有办法很公平地对比二者的后端, 目前学术论文主要是在 LLVM 内实现一个特定的 allocator 和 production allocator 进行对比.
+
+有趣的是, 连 GCC IRA 的主要作者 Vladimir Makarov 在 Reddit 上讨论 LLVM 和 GCC RA 时, 更多也是给经验判断. 他指出真实 RA 的表现受到 splitting, rematerialization, target constraints, coalescing 等大量细节影响, 他个人经验是 graph-coloring 在复杂 CFG 上可能表现更好, 并举 SPEC perl 一类程序作为例子.
+
+另外, 根据我找到的一些 benchmark 来看, GCC 表现普遍强于 LLVM, 不论是编译时间还是代码生成的质量, 而我看到知乎上很多人说 LLVM 明显比 GCC 快, 我觉得这一点有失偏颇. 在大型 C++ 项目中链接的速度也很重要, 而 LLVM 的 LLD 远强于 GNU 的 ld, 论据如下:
+
+| target         |   GNU ld |     LLD |
+| -------------- | -------: | ------: |
+| ffmpeg debug   |   1.72 s |  0.35 s |
+| mysqld debug   |   8.50 s |  0.68 s |
+| clang debug    | 104.03 s |  5.28 s |
+| chromium debug | 209.05 s | 16.70 s |
+而且还有一个很有意思的证据, mold 又进一步把 LLD 拉开了.
+
+mold 当前公布的大型程序 benchmark 是:
+
+| Program      | GNU ld | GNU gold |   LLD |      mold |
+| ------------ | -----: | -------: | ----: | --------: |
+| MySQL 8.3    | 10.84s |    7.47s | 1.64s | **0.46s** |
+| Clang 19     | 42.07s |   33.13s | 5.20s | **1.35s** |
+| Chromium 124 |      — |   27.40s | 6.10s | **1.52s** |
+所以我觉得没必要批评 GCC 已经被 LLVM 击败了, 真正导致编译慢的是 GNU 的 Binutils, 它要背大锅. 一个比较值得尝试的方向是 GCC + LLD 或者 GCC + mold.
diff --git a/content/notes/register-allocation/ssa/index.md b/content/notes/register-allocation/ssa/index.md
new file mode 100644
index 0000000..443f2b1
--- /dev/null
+++ b/content/notes/register-allocation/ssa/index.md
@@ -0,0 +1,443 @@
+---
+title: "第七部分: SSA 与寄存器分配"
+slug: "ssa"
+lang: zh-Hans
+series: register-allocation
+weight: 80
+created: 2026-08-17
+updated: 2026-09-09
+license: CC-BY-SA-4.0
+---
+
+前面几部分讨论寄存器分配时, 我们一直从一般的 live range 出发. 一个机器级虚拟寄存器可能经历多次 definition, allocator 再通过 liveness 判断某一个 definition 产生的值在什么区域有效. SSA, 即 Static Single Assignment, 把这个关系整理得更规整: 每个 SSA value 只有一个 definition, 每个 use 都明确引用这个 definition 产生的值. 这种结构首先简化了 def-use 关系, 随后又影响 live range 的形状, interference graph 的性质, phi 的处理方式以及 SSA destruction.
+
+SSA 对寄存器分配的意义不只是 "变量改了名字". 对一个 allocator 来说, single assignment 带来的真正收益是每个 value 都有唯一的生命周期起点. 如果再结合 dominance, 很多在一般 live range 上需要额外分析的关系会获得比较严格的结构.
+
+## 1 Single Definition, Dominance 与 Live Range
+
+先看普通命令式形式:
+
+```text
+x = a + b
+use x
+
+x = c + d
+use x
+```
+
+如果把 `x` 当作同一个名字, 它在这里显然有两个不同 definition. SSA 会把它们拆成两个 values:
+
+```text
+x1 = a + b
+use x1
+
+x2 = c + d
+use x2
+```
+
+现在 `x1` 的所有 uses 都属于第一个 definition, `x2` 的所有 uses 都属于第二个 definition. 对 liveness 来说, allocator 不再需要询问 "这里使用的是 `x` 的哪一次定义"; SSA name 已经把这个问题编码进 IR.
+
+SSA 还要求一个 value 的 definition dominate 它的普通 uses. 假设 `v` 在 basic block `D` 中定义, 并在 basic block `U` 中使用, 那么从函数入口到达这个 use 的每一条控制流路径都必须经过 `D`. 例如:
+
+```text
+        B1
+        |
+      v = ...
+       / \
+      /   \
+     B2   B3
+     |     |
+   use v   |
+      \   /
+       \ /
+        B4
+        |
+      use v
+```
+
+这里 `B1` 位于到达两个 uses 的所有控制流路径上. 因此 `v` 的生命周期有一个明确的根, 就是 `B1` 中的唯一 definition.
+
+如果从 dominance tree 上观察一个 SSA value 的 live region, 它通常可以表示成一棵连通子树. 一个 use 要求 value 从 definition 一直沿相应的 dominance 路径保持有效, 多个 uses 对应的这些路径合起来仍然保持连通. 这条性质后面会直接导出 SSA interference graph 的 chordal structure.
+
+需要注意 source variable 和 SSA value 的区别. 一个源语言变量可能在 SSA 中变成很多 names, 每个 name 都是独立 value. Register allocator 更关心这些 values 的 live ranges, 而不会试图让所有来自同一个源变量的 SSA names 永远使用同一个物理寄存器.
+
+## 2 Phi 的 Use 发生在 CFG Edge 上
+
+SSA 在控制流汇合处使用 phi. 例如:
+
+```text
+B1:
+    x1 = ...
+    goto B3
+
+B2:
+    x2 = ...
+    goto B3
+
+B3:
+    x3 = phi(x1 from B1, x2 from B2)
+    use x3
+```
+
+Phi 的 operand 和普通指令 operand 有一个根本区别. `x1` 的 use 属于 edge `B1 -> B3`, `x2` 的 use 属于 edge `B2 -> B3`. 程序从 `B1` 进入 `B3` 时选择 `x1`, 从 `B2` 进入时选择 `x2`.
+
+```text
+B1 ---- x1 ----\
+                \
+                 >---- B3: x3 = phi(...)
+                /
+B2 ---- x2 ----/
+```
+
+因此不能把 `x1` 和 `x2` 简单地都加入 `LIVE_IN[B3]`. 这样做会让两个来自互斥 predecessor edges 的 operands 在 `B3` 入口处看起来同时 live, 进而制造虚假的 interference.
+
+Block-level liveness 在处理 phi 时通常需要保留 edge-specific 信息. 假设 successor `S` 中有一个 phi, predecessor `B` 只应该把属于 `B -> S` 这条 edge 的 phi operand 加入自己的 live-out. 来自其他 predecessors 的 operands 与 `B` 无关.
+
+Phi result `x3` 则在 `B3` 入口处产生. 所以从数据流的角度看, 可以把 phi 想象成多条 predecessor edges 上发生的 value transfer, 最终在 successor 入口得到一个新的 SSA value.
+
+这种语义对寄存器分配很有利. 如果 `x1`, `x2`, `x3` 没有其他 interference 阻止它们共享寄存器, allocator 可以安排 $x_1 \mapsto R_0$, $x_2 \mapsto R_0$, $x_3 \mapsto R_0$. SSA destruction 以后, 两条 edge 上都不需要实际的数据移动.
+
+如果得到 $x_1 \mapsto R_0$, $x_2 \mapsto R_1$, $x_3 \mapsto R_0$, 那么只有 `B2 -> B3` 需要一次从 `R1` 到 `R0` 的 copy. 因此 phi 会自然形成 coalescing preference: incoming value 和 phi result 倾向于使用相同的物理寄存器.
+
+这里同样不能只根据 phi 自身判断两个 operands 是否 interfere. `x1` 和 `x2` 可能因为程序其他位置的 liveness 而发生冲突, 也可能完全没有冲突. Phi 的 edge semantics 只是保证 allocator 不会因为 phi 这一条指令本身凭空制造冲突.
+
+## 3 SSA Interference Graph 为什么是 Chordal Graph
+
+SSA 给 graph coloring 带来的理论优势, 来自前面提到的 dominance structure.
+
+把 dominance tree 看成一棵真正的树. 每个 SSA value 的 live region可以对应其中的一棵连通子树. 两个 values 发生 interference, 意味着它们的 live regions 在某个程序位置相交. 因此 SSA interference graph 可以理解为一组 dominance-tree subtrees 的 intersection graph.
+
+例如 dominance tree 为:
+
+```text
+          A
+        /   \
+       B     C
+      / \     \
+     D   E     F
+```
+
+假设几个 values 的 live regions 大致为:
+
+```text
+v1: A - B - D
+
+v2:     B - D - E
+
+v3:             C - F
+```
+
+`v1` 和 `v2` 的 live regions 相交, 因此在 interference graph 中存在边. `v3` 所在的区域与前两者分离, 因而这组 live regions 本身不会生成 `v1-v3` 或 `v2-v3` 的 interference.
+
+树的连通子树所形成的 intersection graph 是 chordal graph. 因此, 在保持 strict SSA liveness 结构的条件下, SSA interference graph 具有 chordal property.
+
+Chordal graph 中, 所有长度至少为 4 的无弦环都会被排除. 例如:
+
+```text
+a ----- b
+|       |
+|       |
+d ----- c
+```
+
+这是一个没有 chord 的四边形. 如果加入 `a-c`:
+
+```text
+a ----- b
+|  \    |
+|    \  |
+d ----- c
+```
+
+环上出现了一条连接非相邻顶点的 chord.
+
+Chordal graph 存在 perfect elimination ordering. 按这种顺序逐个删除节点时, 每个被删除节点当前的邻居都会形成 clique. 反过来按照这个顺序进行 greedy coloring, 可以得到最优 coloring. 因而对于 chordal graph 有:
+
+$$  
+\chi(G)=\omega(G)  
+$$
+
+这里 $\chi(G)$ 是 chromatic number, $\omega(G)$ 是最大 clique 大小.
+
+这个等式和第三, 四部分讨论的一般 interference graph 差别很大. 对一般图, $\omega(G)$ 只是 $\chi(G)$ 的下界. 最大 clique 只需要 3 种颜色, 整张图仍然可能需要 4 种甚至更多颜色. 对 chordal graph, 最大 clique 已经决定了最少颜色数.
+
+如果当前 register class 有 $K$ 个完全等价的物理寄存器, 并且 SSA interference graph 满足 $\omega(G)\le K$, 那么单纯从 graph coloring 的角度看, 存在合法的 $K$-coloring, 而且可以高效找到.
+
+这里的限制条件需要保留. 真实 RA 还有 register classes, pre-colored nodes, physical register aliasing, fixed constraints 和 register tuples. 这些约束加入以后, "最大 clique 不超过 $K$" 已经不能单独保证最终 machine assignment 成功. Chordal property 解决的是基础 interference coloring 问题.
+
+## 4 SSA 把很多困难推向了 Spilling
+
+假设一个 register class 只有 $K$ 个物理寄存器, 某个位置却有 $K+2$ 个 SSA values 同时 live. 这些 values 在 interference graph 中形成至少大小为 $K+2$ 的 clique. 即使 graph 是 chordal, coloring 也不可能把 $K+2$ 个互相冲突的 values 塞进 $K$ 个寄存器.
+
+所以在 SSA-based register allocation 中, spilling 和 coloring 可以形成比较清晰的分工. Spilling 或 splitting 负责降低某些位置的 register pressure, 随后的 coloring 再利用 SSA interference graph 的结构完成 assignment.
+
+这一点和 Chaitin-Briggs 的组织方式有所区别. Chaitin-style algorithm 在 simplify 过程中遇到 high-degree structure 时选择 potential spill, coloring 和 spill decision 彼此交织. SSA-based allocator可以先研究哪些 values 需要从高压力区域移走, 在压力满足要求以后再利用 chordal graph 进行 coloring.
+
+SSA 也方便 spill rewrite 保持 single-definition property. 假设:
+
+```text
+v1 = ...
+...
+use v1
+...
+use v1
+```
+
+如果 `v1` 被 spill, 可以写成:
+
+```text
+v1 = ...
+store [slot], v1
+
+...
+
+t1 = load [slot]
+use t1
+
+...
+
+t2 = load [slot]
+use t2
+```
+
+`t1` 和 `t2` 各自只有一个 definition, 因而仍然可以作为新的 SSA values. 原来的长 live range 被分解成多个短 live ranges, allocator 可以继续对这些新 values 使用 SSA liveness.
+
+Reload 放在哪里会影响新的 register pressure. 假设一个 spilled value 在两个分支中都有 use:
+
+```text
+        B1
+       /  \
+      B2  B3
+      |    |
+    use   use
+```
+
+一种方案是在 `B1` 中 reload 一次:
+
+```text
+        load t
+          |
+         / \
+        B2 B3
+        |   |
+      use  use
+```
+
+这样两个分支共享一个 reload, 但 `t` 从 `B1` 开始就保持 live, 生命周期比较长. 另一种方案是在 `B2` 和 `B3` 中分别 reload, 产生两个较短的 SSA values. 后者可能增加动态 memory operations, 同时降低寄存器压力.
+
+Spill placement 因而同时涉及 execution frequency 和 live range length. SSA 使新 definitions 和 uses 的关系更清楚, 却没有消除这种 cost trade-off.
+
+## 5 Live Range Splitting 在 SSA 中如何表示
+
+Live range splitting 在 SSA 里可以理解成主动创建新的 definitions.
+
+假设原来:
+
+```text
+v1:
+|-----------------------------------|
+```
+
+allocator 希望把它拆成三个 fragments:
+
+```text
+v1:
+|-----------|
+
+v2:
+             |-----------|
+
+v3:
+                          |--------|
+```
+
+那么 `v2` 和 `v3` 可以成为新的 SSA values. 后续 uses 根据程序位置重新命名, 分别引用对应 fragment. Fragments 之间通过 copy, spill/reload 或 phi 传递值.
+
+在直线代码里可能得到:
+
+```text
+v1 = ...
+...
+v2 = COPY v1
+...
+v3 = COPY v2
+...
+```
+
+如果最后三个 fragments 都分配到同一个 physical register, 这些 copies 可以被消除. 如果中间 fragment需要进入 memory, transfer 就可能转化成 store 和 reload.
+
+带控制流的 splitting 会更接近 SSA construction. 例如两个分支分别产生新的 fragments, 在汇合以后需要继续使用同一个逻辑值:
+
+```text
+          v1
+         /  \
+        /    \
+      v2      v3
+        \    /
+         \  /
+      v4 = phi(v2, v3)
+```
+
+这里 `v2` 和 `v3` 分别是不同路径上的 definitions, `v4` 在 join block 中重新把它们合并成一个 SSA value. 所以 SSA-based splitting 往往需要 dominance, renaming 和 phi placement 协同工作.
+
+从 allocator 的角度看, splitting 的收益仍然是缩短 allocation object. 原来一个 value 跨越多个高压力区域, 拆开以后每个 fragment 可以分别选择 physical register 或 memory location, 不需要整个生命周期保持同一种 allocation decision.
+
+## 6 Coalescing 与 SSA Structure 的关系
+
+Phi destruction 会产生 copy, allocator 又希望通过 coalescing 删除这些 copies. 但 coalescing 会改变 SSA value 原本整齐的 live-range structure.
+
+考虑:
+
+```text
+x3 = phi(x1, x2)
+```
+
+`x1`, `x2`, `x3` 分别拥有自己的唯一 definition. 如果 allocator 把它们全部合并为一个 allocation object, 这个联合对象会包含多个 definitions. 从物理寄存器分配角度看这完全可能是理想结果, 因为三个 values 可以共享一个 physical register; 从 SSA structural analysis 的角度看, 联合对象已经不再对应一个单一 definition 支配的 live subtree.
+
+任意识别两个不相邻节点进行 coalescing, 还可能破坏原来的 chordal interference structure. 因此某些 SSA-based allocators 会尽量在利用完 SSA coloring 性质之后再进行更激进的 coalescing, 或者只允许满足特定条件的合并.
+
+这里和第五部分的 IRC 有相似的问题背景. IRC 从一般 interference graph 出发, 使用 Briggs 或 George criterion 控制 coalescing 风险. SSA-based RA 拥有更规整的初始图, 因而更加有理由保护这份结构, 至少在需要利用 chordal coloring 的阶段不要随意破坏它.
+
+Phi 本身仍然提供了很强的 coalescing preference. 因此实际设计需要在两件事之间选择时机: 一方面希望保留 SSA structure 方便 allocation, 另一方面希望尽量让 phi related values 获得同一 physical register.
+
+## 7 SSA Destruction 与 Parallel Copy
+
+最终机器代码不能保留抽象 phi, 因此在某个阶段需要进行 SSA destruction.
+
+对于:
+
+```text
+B1:
+    ...
+    goto B3
+
+B2:
+    ...
+    goto B3
+
+B3:
+    x3 = phi(x1 from B1, x2 from B2)
+```
+
+可以把它展开成两条 edge-specific transfers:
+
+```text
+B1 -> B3:
+    x3 <- x1
+
+B2 -> B3:
+    x3 <- x2
+```
+
+如果 `x1` 和 `x3` 已经分配到同一寄存器, 第一条 transfer 无需生成真正的 move. `x2` 和 `x3` 同理.
+
+一个 block 同时存在多个 phi 时, 某条 predecessor edge 上会出现一组同时发生的 assignments:
+
+```text
+a3 <- a1
+b3 <- b1
+c3 <- c1
+```
+
+这里需要使用 parallel copy semantics. 所有右侧值都按照 assignment 开始前的状态读取.
+
+假设物理寄存器分配以后出现:
+
+```text
+R0 <- R1
+R1 <- R0
+```
+
+如果先执行 `R0 <- R1`, 原来 `R0` 中的值已经丢失. 因此需要 temporary:
+
+```text
+TMP <- R0
+R0  <- R1
+R1  <- TMP
+```
+
+对于没有 cycle 的 copy dependencies, compiler 可以安排一个安全顺序. 例如:
+
+```text
+R0 <- R1
+R2 <- R0
+```
+
+如果两个 source 都表示旧寄存器值, 应该先执行:
+
+```text
+R2 <- R0
+R0 <- R1
+```
+
+这样旧 `R0` 在被覆盖之前已经保存到 `R2`.
+
+因此 SSA destruction 通常先把 phi 转换成 edge-specific parallel copies, 再由 parallel-copy resolver 把它们变成串行机器指令. Coalescing 做得越好, 真正需要 resolver 处理的 copies 就越少.
+
+## 8 Critical Edge 与 Copy Placement
+
+Phi 的 transfers 属于 CFG edges, 所以 copy placement 还会遇到 critical edge.
+
+假设:
+
+```text
+        B
+       / \
+      v   v
+      S   X
+      ^
+      |
+      P
+```
+
+`B` 有两个 successors, `S` 也有多个 predecessors, 因而 `B -> S` 是 critical edge. 如果某条 phi transfer 只属于 `B -> S`, 那么不能直接放在 `B` 末尾, 因为走 `B -> X` 时也会执行; 放到 `S` 开头也会影响从 `P -> S` 进入的路径.
+
+常见做法是 split 这条 edge:
+
+```text
+        B
+       / \
+      v   v
+      E   X
+      |
+      v
+      S
+      ^
+      |
+      P
+```
+
+新的 block `E` 只会在原来的 `B -> S` 路径上执行, 因而可以安全放置对应 copies.
+
+这会说明 SSA destruction 在 compiler pipeline 中并不是一个纯粹的语法转换. 它可能修改 CFG, 引入 copies, 改变 liveness, 进而影响寄存器分配. 如果很早 destruction, 后面的 allocator 要处理更多 ordinary copies; 如果延迟 destruction, allocator 本身就需要理解 phi operands 的 edge semantics.
+
+## 9 Machine SSA 与工业寄存器分配
+
+SSA-based RA 并不等于工业编译器一定使用某一种专门的 "SSA coloring algorithm". SSA 也可以只是寄存器分配之前维持的一种 machine-level invariant.
+
+LLVM 就很适合说明这一点. Instruction selection 以后, Machine IR 中的 virtual registers 在相当一段 pipeline 内仍然保持 SSA 性质. 一个 virtual register 通常具有唯一 definition, machine passes 可以沿 def-use chains 查询它的数据来源和使用位置.
+
+寄存器分配逐渐改变这种状态. Coalescing 可能把多个 values 连接起来, splitting 会生成新的 fragments, spilling 会插入 stores 和 reloads, 最终 virtual registers 被改写成 physical registers. 一个 physical register 如 `RAX` 在整个函数中会反复被不同指令定义, 自然不再满足 SSA.
+
+所以工业 allocator 经常会经历这样的变化:
+
+```text
+Machine SSA
+    |
+    | virtual registers with single defs
+    v
+Coalescing / splitting / spilling
+    |
+    v
+General live ranges
+    |
+    v
+Physical register assignment
+    |
+    v
+Non-SSA physical-register code
+```
+
+LLVM 的 `LiveIntervals` 之所以需要 `VNInfo` 一类结构, 也与这种变化有关. 一旦一个 register live range 中出现多个 definitions, allocator 就需要区分不同 definitions 产生的 value numbers, 单纯依赖最初的 SSA name 已经不够.
+
+SSA 给寄存器分配提供了一块结构比较规整的起点. Definition 和 use 的关系清楚, phi operands 具有 edge semantics, live regions 与 dominance tree 联系紧密, interference graph 也因此具有 chordal structure. 进入下一部分以后, 我们会开始把这里相对抽象的 "颜色" 逐步替换成真实 CPU 的寄存器资源, 讨论 register class, calling convention, pre-colored register, fixed operand, subregister aliasing, two-address constraint 和 register pair 等机器约束.