TPU 调度 Pipeline

地址适用于 libtpu-0.0.40-cp314 wheel 中的 libtpu.so（build-id 89edbbe81c5b328a958fe628a9f2207d）。其他版本会有所不同。

摘要

TPU 代码会被调度两次，发生在两个不同 IR 层级，使用两个不同算法，并且第二次调度依赖一个独立的资源分配 pass；这一点在 CPU 或 GPU 后端中没有对应物。本页是这个栈的地图。第一个调度器是 HLO 层级的 LatencyHidingScheduler（LHS），一个 greedy、bottom-up、critical-path list scheduler；它会重新排序已经做过内存最小化的 HLO 序列，使长延迟的 async 工作（collective start、host/ICI DMA、async copy）尽早发出，并隐藏在独立计算之下。第二个是 LLO 层级的 bundle packer，一个前向 greedy 的 "earliest-legal-bundle" list scheduler；它把每个低层 op 分配到固定宽度 VLIW bundle 内的 slot，再加上一个对内部硬件循环做 software pipeline 的独立 modulo scheduler。二者之间是 MXU-sequence assignment 和 MRB（matrix-result-buffer）allocation pass，它决定 matmul 累加链和 result-FIFO placement，而 bundle packer 的 latch ordering 随后依赖这些结果。为所有这些决策定价的是第 VII 部分 cost model 中的 ResourceType 模型和 AsyncTracker。

熟悉 LLVM 的读者可以逐阶段类比。LHS 是经典 MachineScheduler 风格的 ready-set list scheduler，但它反向构建序列（current_time 时钟从 root 向回走），并且其优先级主要由 async critical-path depth/height 主导，而不是普通 dependency height。bundle packer 相当于 TPU 版 LLVM VLIWPacketizer（DFAPacketizer），只是资源模型是逐代 BundleRequirement bitmap，而不是 deterministic finite automaton；并且没有 spill-to-next-bundle 搜索，slot 满时会追加一个新的空 bundle，然后重试该 op。modulo scheduler 是教科书式 iterative modulo scheduling：基于 LLVM ScheduleDAGMI over MachineInstr 做 RecMII/ResMII initiation-interval 搜索。MXU/MRB pass 则完全没有 LLVM 对应物：它是面向脉动阵列累加 buffer 的领域特定 bin-packer。

这是一个导览页。它框定四个阶段，固定阶段顺序，并链接到记录细节的各个子页。它不重复算法本身；LHS comparator 的 22 个 key、逐代 bundle slot matrix、MRB reservation timeline、modulo II 搜索，都分别位于下方交叉引用的独立页面。

对重新实现而言，这个栈的契约是：

两个调度器操作互不相交的 IR，且从不共享状态。 LHS 重排 HloInstruction 并写入 module->schedule；bundle packer 把 LloInstruction（或 MachineInstr）重排/打包进 Bundle 记录。HLO 调度 annotation 不会流入 bundle packing；它们只引导 HLO computation 的 async-overlap。
LHS 要求已有 base schedule 作为前置条件。 它不是从零构建 schedule；HloMemorySchedulerWithBrkgaFallback 会先放下一个最小化 memory-pressure 的顺序，LHS 再用一个能重叠 async edge 的顺序替换它。四个 LHS 页面是同一个共享主体的变体；它们只在 pipeline placement 和 SchedulerConfig/AsyncTracker 输入上不同。
MXU-sequence assignment 运行在两个调度器之间，并向 packer 供给信息。 AssignMxusForSequenceGroup 决定 matmul 累加链、latch placement 和 MRB FIFO entry；bundle packer 的逐 op latch ordering，以及逐代 encoder 的 latch serialization，会在下游读取这个 assignment。
每个调度决策都由 bundle cost model 定价。 Node cost（对 23-slot bundle vector 执行 MaxResourceCycles）和 dependency latency（LatencyBetween）供给 LHS 优先级；ResourceType/AsyncTracker 模型 gate 哪些 async op 可以按物理资源重叠。


Stage 1 — HLO scheduler	`xla::LatencyHidingScheduler::RunImpl` @ `0x136321a0`；per-comp drain `DefaultSchedulerCore::ScheduleComputation` @ `0x1362eb60`；comparator `FindAndExtractBestNodeAvailable` @ `0x13618880`
Stage 1 driver	`(anon)::RunHloScheduler` @ `0x1096fac0` — two pipelines：`final_scheduler`（base）→ `async_scheduling`（LHS）
Stage 2 — MXU/MRB assignment	`AssignMxusForSequenceGroup` @ `0x10f753c0`（worker `0x10f77ca0`）；`MrbChainAllocator::ExtendMrbReservation` @ `0x10f58800`；`MxuAssigner::AllocateMrbEntriesAsFifo` @ `0x10f3ef80`，`BounceBetweenMsrs` @ `0x10f3fae0`
Stage 3 — LLO bundle packer	`PackBundles` @ `0x10a30a20`；`GlobalBundlePacker::Pack` @ `0x10a86420`；`BundlePacker::Feed` @ `0x14021f20`
Stage 3b — modulo（loops）	`llvm::TPUScheduleDAGModulo::findSchedule` @ `0x13b1d7c0`；II search `calculateResourceMII` @ `0x13c0bee0`，`calculateLargestLatencyMII` @ `0x13c0b840`
定价 — cost model	`MaxResourceCycles` @ `0x1c89b9e0`；`LatencyBetween` @ `0x1c89f820`；`TpuAsyncTracker::GetResourceHazardType` @ `0x110015e0`
IR 层级	Stage 1：HLO（`HloInstruction`）· Stages 2–3：LLO（`LloInstruction`）/ LLVM `MachineInstr`
置信度	已确认（byte-anchored），除非某行或 callout 另有说明

阶段顺序

四个阶段严格按 pipeline 顺序运行，每个阶段消费前一阶段的输出。两个 list scheduler 位于 matrix-resource assignment 的两侧；cost model 会被三者查询。

text

HLO module (post layout-assignment, post main-fusion)
   │
   ▼
┌──────────────────────────────────────────────────────────────────────┐
│ STAGE 1 — HLO LATENCY-HIDING SCHEDULER          RunHloScheduler        │
│   final_scheduler   : HloMemorySchedulerWithBrkgaFallback (base order) │   0x1096fac0
│   async_scheduling  : LatencyHidingScheduler::RunImpl (overlap rewrite)│   0x136321a0
│     └─ DefaultSchedulerCore::ScheduleComputation (reverse list drain)  │   0x1362eb60
│          └─ FindAndExtractBestNodeAvailable (22-key ReadySetLt)        │   0x13618880
│     priced by  GetLatencyBetween → MaxResourceCycles / LatencyBetween  │
│     gated by   TpuAsyncTracker (ResourceType hazard classes)           │
└──────────────────────────────────────────────────────────────────────┘
   │  (HLO → MHLO/TLP → tpu dialect → LLO IR descent; not on this page)
   ▼
┌──────────────────────────────────────────────────────────────────────┐
│ STAGE 2 — MXU SEQUENCE + MRB ALLOCATION                                │
│   AssignMxusForSequenceGroup (bin-pack matmuls into MXU sequences)     │   0x10f753c0
│     └─ MrbChainAllocator (accumulation-chain reservation timeline)     │   0x10f58800
│     └─ AllocateMrbEntriesAsFifo / BounceBetweenMsrs (result FIFO/MSR)  │   0x10f3ef80
│     └─ SetLatchIndices (commit latch ids onto each MxuSequence)        │   0x10f3b4c0
└──────────────────────────────────────────────────────────────────────┘
   │  (latch indices + MRB entries are inputs to slot legality below)
   ▼
┌──────────────────────────────────────────────────────────────────────┐
│ STAGE 3 — LLO BUNDLE PACKING                    PackBundles            │   0x10a30a20
│   GlobalBundlePacker::Pack (forward greedy, topological)               │   0x10a86420
│     └─ BundlePacker::Feed (earliest-legal-bundle + SlotTracker)        │   0x14021f20
│     └─ per-gen TpuBundleRestrictions (slot legality, latch ordering)   │
│   STAGE 3b (inner loops only): TPUScheduleDAGModulo (II-search SWP)    │   0x13b1d7c0
└──────────────────────────────────────────────────────────────────────┘
   │
   ▼
vector<Bundle> ─→ per-gen Encoder*::EncodeBundle (latch serialization) ─→ raw bytes
```text

这个顺序不是任意的。Stage 1 在 HLO 上运行，正是因为 async overlap 是一个*宏观*决策，也就是哪个 collective 在哪个 fusion 之前 start；这种决策必须在程序仍是带 whole-tensor dependency edge 的高层 dataflow graph 时做出。Stage 3 在 LLO 上运行，因为 slot legality 是一个*微观*决策，也就是每个具体 op 占用一个 VLIW word 中的哪个 scalar/vector/MXU lane；只有在程序降低到硬件 op 之后，这个问题才存在。Stage 2 位于二者之间，因为 MXU 的 accumulation buffer（MRB）和 result latch 是一种物理资源，其 assignment 比 HLO 更细粒度，但它是 slot legality 的*前置条件*，而不是 slot legality 的结果：bundle packer 需要先知道每个 matmul result 落在哪个 latch，才能决定两个 matmul op 在同一个 bundle 中是否冲突。

> **陷阱 — 存在两个无关的“调度器”，混淆它们会导致所有定价错误。** Stage 1（`LatencyHidingScheduler`）和 Stage 3（`BundlePacker`）都是 list scheduler，但它们不共享代码、不共享 IR，也不共享 cost unit。Stage 1 通过 `MaxResourceCycles` 以 *bundle cycle* 定价，用于重叠 async transfer；Stage 3 通过逐代 `BundleRequirement` bitmap 以 *slot occupancy* 定价，用于装入 VLIW word。只运行其中一个而跳过另一个的重新实现，要么产生没有 async overlap 的 schedule（跳过 Stage 1），要么产生带 slot conflict、无法编码的指令流（跳过 Stage 3）。

---

## Stage 1 — HLO Latency-Hiding Scheduler

### 用途

LHS 是生产用 HLO scheduler。它由 `RunHloScheduler`（`0x1096fac0`）调用，位于 `DeepseaCompilerBase::RunHloPasses` 的 final-scheduling 阶段，在 layout assignment 和 main fusion 之后。它不是从空白开始构建 schedule：`RunHloScheduler` 会先运行一个 `final_scheduler` pipeline，其中 `HloMemorySchedulerWithBrkgaFallback` 放下最小化 memory-pressure 的 base order（LHS 的 `has_schedule()` 前置条件）；然后运行 `async_scheduling` pipeline，其中 `LatencyHidingScheduler` 在 TPU resource model 下，用一个把 async start/done pair 与 compute 重叠的顺序替换该 base order。

### 结构

```text
RunHloScheduler                                   0x1096fac0
  ├─ pipeline "final_scheduler"
  │     └─ HloMemorySchedulerWithBrkgaFallback     ── base memory-minimized order
  └─ pipeline "async_scheduling"
        └─ AddPass<LatencyHidingScheduler>          ── shared RunImpl 0x136321a0
              ├─ SchedulerConfig (per-variant POD)
              ├─ SchedulingContext (LatencyEstimator + TpuAsyncTracker)
              └─ DefaultSchedulerCore::ScheduleComputation  0x1362eb60
                    └─ FindAndExtractBestNodeAvailable       0x13618880

最重要的结构事实是，RunImpl 主体（0x136321a0）在所有变体之间共享。编译出的 RunImpl 只有一个，也就是未修改的上游 latency_hiding_scheduler.cc 主体。下面四个 LHS 页面记录的变体，只在交给 AddPass<LatencyHidingScheduler> 的两个输入上不同：SchedulerConfig POD 和 SchedulingContext（estimator + async tracker）。核心算法，包括反向 drain loop、22-key ReadySetLt comparator、预计算 async depth/height、memory-pressure retry loop，只在 core page 上记录一次。

变体映射

变体	改动内容	页面
Core	共享的 `RunImpl` 主体：drain loop、comparator、async tracker、memory retry	LatencyHidingScheduler Core
post_layout	0.0.40 中的实际 placement：two-pipeline `RunHloScheduler`，`SchedulerConfig`/estimator/tracker wiring	LHS post_layout
post_layout_pre_fusion	同一主体接到一个 dead "Pre main fusion" slot；只有 config delta	LHS post_layout_pre_fusion
ILP	`EnableIlpLatencyHidingScheduler` gate 只替换 async classifier；comparator 和 `RunImpl` 不变。具名 `ILPMemoryScheduler` MIP 是独立的 opt-in memory scheduler（可通过 `GetMemorySchedulerAlgorithm` case 6 到达），不是由该 flag 调用	LHS ILP variant

注意 — "ILP-LHS" flag 并不会引入 ILP scheduler。 flag family xla_tpu_enable_ilp_latency_hiding_scheduler 的名字看起来像是一个 integer-linear-programming scheduler 替换了 greedy list scheduler。事实并非如此：RunHloScheduler 内的 EnableIlpLatencyHidingScheduler 路径只替换哪些 async-op kind 会作为 overlap candidate 暴露给同一个 greedy LHS（comparator 和 RunImpl 不变）。这不同于真正的 xla::ILPMemoryScheduler CP-SAT memory scheduler；后者是一个独立、可达的 opt-in：GetMemorySchedulerAlgorithm（0x10abd6a0）在其 0..9 switch（Default0 / List1 / DFS2 / PostOrder3 / BRKGA4 / BFS5 / ILP6 / Backtracking7 / BruteForce8 / LocalOrder9）的 case 6（MemorySchedulerProto.Value::ILP）选择它，key 来自 TpuCompilationEnvironment+0x10c。该 MIP 是按需的（不是默认），不是 dead。见 ILP variant page。

消费内容

comparator 的 per-node cost 和 per-edge latency 来自 cost model：CostModelLatencyEstimator::GetLatencyBetween（0x10ff8f00）把 compute edge 路由到 CostModel::GetCycles → MaxResourceCycles（bundle throughput），把 dependency edge 路由到 LatencyBetween（RAW stall）。两个 async op 是否可以重叠，则由 TpuAsyncTracker resource model 决定：这是一个 0..46 的 ResourceType enum，带有 per-resource hazard class。cost feed 和 resource model 各有页面：Bundle-Aware Cost 和 ResourceType Taxonomy。

Stage 2 — MXU Sequence Assignment 与 MRB Allocation

用途

在 HLO scheduler 和 bundle packer 之间，一个 TPU 专用 pass 会把 matmul op 分配到脉动阵列的物理资源上。每条 matmul 累加链成为一个 MxuSequence；这些链被装箱到可用 MXU pass 上（AssignMxusForSequenceGroup，0x10f753c0，worker 0x10f77ca0）；保存 partial sum 的 matrix-result buffer（MRB）会在时间线上被预留（MrbChainAllocator，0x10f58800）；排空这些 buffer 的 result FIFO / matrix-shift-register（MSR）会被放置（AllocateMrbEntriesAsFifo 0x10f3ef80，BounceBetweenMsrs 0x10f3fae0）。assignment 通过把 latch index 写到每个序列上来 commit（SetLatchIndices，0x10f3b4c0）。

这个阶段在 scalar backend 中没有对应物。它存在是因为 TPU MXU 是一个 reservation-scheduled 单元：matmul 不是一个周期完成，而是把 weight 和 operand 推入 systolic pipeline，其结果很多周期后出现在特定 buffer 中；compiler 必须静态预留该 buffer 以及读取它的 latch。下游 bundle packer 把这些 latch index 作为 slot-legality 输入，也就是两个会撞到同一个 latch 的 op 不能共享一个 bundle。

结构

text

AssignMxusForSequenceGroup                        0x10f753c0   ── bin-pack matmul sequences onto MXU passes
  ├─ (worker) AssignMxusForSequenceGroupInternal  0x10f77ca0
  ├─ MrbChainAllocator                                          ── reservation timeline for accumulation buffers
  │     ├─ ExtendMrbReservation                   0x10f58800
  │     ├─ SplitAccumulationChain                 0x10f598e0
  │     ├─ ReleaseMrbReservation                  0x10f5f9e0
  │     └─ AdvanceTimeTo                          0x10f5e9e0
  ├─ AllocateMrbEntriesAsFifo                     0x10f3ef80   ── result FIFO placement
  ├─ BounceBetweenMsrs                            0x10f3fae0   ── matrix-shift-register ping-pong
  └─ SetLatchIndices                              0x10f3b4c0   ── commit latch ids onto each MxuSequence
```text

### 子页

| 组件 | 负责内容 | 页面 |
|---|---|---|
| `MxuSequence` / `SequenceInfo` | 每序列记录和 `set_mxu` commit | [MxuSequence / SequenceInfo](mxu-sequence-struct.md) |
| `AssignMxusForSequenceGroup` | 把 sequence 放到 MXU pass 上的 bin-packing 算法 | [MXU Assignment Bin-Packer](mxu-assignment-binpacker.md) |
| `MrbChainAllocator` | accumulation-chain reservation timeline 与 jitter model | [MRB Chain Allocator](mrb-chain-allocator.md) |
| `AllocateMrbEntriesAsFifo` / `BounceBetweenMsrs` | result-FIFO 与 MSR placement | [MRB FIFO / MSR Placement](mrb-fifo-msr-placement.md) |
| `SetLatchIndices` / overrun handshake | latch-index assignment 与逐代 overrun handshake | [Latch Assignment & Overrun](latch-assignment-overrun.md) |

> **注意 — Stage 2 是 Stage 3 的依赖，而不是它的 sub-pass。** bundle packer（`BundlePacker::Feed`）在询问某个 op 的 `BundleRequirement` 是否能放进 candidate bundle 时，会读取每个 op 已分配的 latch。把 MXU assignment 折叠进 packer 逐 op 循环的重新实现会造成依赖死锁：latch assignment 需要看到完整 accumulation chain，而 packer 的一次一个 op 前向遍历没有这个视野。

---

## Stage 3 — LLO Bundle Packing

### 用途

bundle packer 接收 LLO instruction list，也就是 raw bundle-byte emission 之前的最后一个 IR，并把每条指令分配到固定宽度 VLIW bundle 内的一个 slot，同时满足逐代资源约束（functional-unit count、vector source port、immediate slot、predicate field）。输出是一个 `vector<Bundle>`，供逐代 `Encoder*::EncodeBundle` 消费。算法是**前向 greedy earliest-legal-bundle list scheduling**：按拓扑顺序遍历 region；对每个 op，先计算其 RAW dependency 允许的最早 bundle，然后请求 `SlotTracker` 找到从该点开始第一个能容纳该 op 的 `BundleRequirement` 的 bundle；若没有可容纳的 bundle，就追加一个新的空 bundle 并重试。没有 spill-to-next-bundle 搜索，也没有 backtracking。

### 两个 Packer，一个算法

libtpu 在不同 IR 层级提供了两个结构相同的 packer：规范的 `xla::jellyfish::BundlePacker` 操作 `LloInstruction*`（直接通往逐代 encoder 的路径），以及一个 `llvm::(anonymous)::BundlePacker` `MachineFunctionPass` 操作 `MachineInstr*`（通过 LLVM MC 的路径）。两者都产生相同的 bundle byte layout；它们是同一个 earliest-legal-bundle 算法在两种 IR 上的实现。逐代 slot legality 封装在四个 `TpuBundleRestrictions` 子类中：`JellyfishBundleRestrictions`、`PufferfishBundleRestrictions`、`ViperfishBundleRestrictions` 和 `GhostliteBundleRestrictions`（二进制中仅存在这几个 `*BundleRestrictions` 子类；没有 `dragonfish`/`6acc60406` restriction table），每个子类都提供 `SetLimits`、`AddXluRequirements`、`AddMxuRequirements` 和 `MatchScalar` virtual。

### 结构

```text
PackBundles                                       0x10a30a20   ── top-level entry (timer, numbering, options)
  └─ GlobalBundlePacker::Pack                      0x10a86420   ── orchestrator, topological region walk
        ├─ PackPhis (region)                       0x10a87160   ── SSA-resolution moves into predecessor tails
        ├─ PackInstruction (instr)                 0x10a875a0   ── regular ops
        │     └─ BundlePacker::Feed (instr, lat)   0x14021f20   ── earliest-legal-bundle + SlotTracker
        │           └─ SlotTracker::FindFeasibleBundleAfter  0x140340a0
        └─ PackBranch (instr)                      0x10a877a0   ── opcodes 0x87/0x88/0xef + delay slots
  → ConsumeBundles                                 0x14026c40   ── hand off vector<Bundle>; ValidatePacking 0x14026ca0

内部循环的 Modulo Scheduling（Stage 3b）

hardware-loop 内部 body 走一条独立路径：LLVM 的 TPUScheduleDAGModulo 执行 iterative modulo scheduling，也就是从 resource 和 recurrence bound 的最大值选择 initiation interval（II），然后把 loop body 放置到 II 个 bundle 上。II 搜索包括 calculateResourceMII（0x13c0bee0，ResMII bound）和 calculateLargestLatencyMII（0x13c0b840，recurrence/RecMII bound）；findSchedule（0x13b1d7c0）是 placement driver。这是教科书式 software-pipelining 路径，只对标记为 hardware_loop 的循环调用；直线 list packer 处理其他所有内容。

子页

组件	负责内容	页面
`GlobalBundlePacker` / `BundlePacker::Feed` / `SlotTracker` / `TpuBundleRestrictions`	前向 greedy 算法、逐代 slot matrix、branch/barrier/phi handling	LLO → Bundle Packing
`TPUScheduleDAGModulo`	hardware loop 的 II-search 与 software-pipelining 路径	Bundle Modulo Scheduling
per-gen `Encoder*::EncodeBundle`	已分配 latch field 如何序列化进逐代 bundle word	Per-Gen Encoder Latch Serialization

怪癖 — packer 从不 spill；它会增长 bundle vector。 当 SlotTracker::FindFeasibleBundleAfter 返回“没有 feasible bundle”时，算法不会搜索替代 slot 或重新排序更早的 op，而是 emplace_back 一个新的空 bundle 并重试同一个 op。只有当单个 op 的 BundleRequirement 即使单独放置也超过该代 per-bundle limit（restriction table 写错）时，编译才会失败；schedule 只是密集并不会导致失败。加入 spill/backtrack heuristic 的重新实现，会在每个 bundle-boundary decision 上偏离二进制。

各阶段如何定价

三个阶段都会查询第 VII 部分 cost model，但粒度和决策不同。

阶段	被定价的决策	Cost query	单位
Stage 1（LHS）	发出哪个 async op、多少 compute 能隐藏其 latency	`GetLatencyBetween` → `MaxResourceCycles`（node bundle cost）+ `LatencyBetween`（edge RAW stall）	bundle cycle → 通过 TC clock 转为秒
Stage 1（LHS）	两个 async op 是否可以重叠	对 0..46 `ResourceType` enum 调用 `TpuAsyncTracker::GetResourceHazardType`	hazard class（shareable / serial / unsharable / nonextendable）
Stage 2（MXU）	MRB reservation window、accumulation jitter	逐代 `CycleTable` matmul/matprep latency	MXU pass cycle
Stage 3（packer）	一个 op 是否能放入 bundle	逐代 `BundleRequirement` slot bitmap（不是 cost model）	slot occupancy

cost model 本身，也就是 23-slot ResourceVector、MaxResourceCycles 的 max-over-lanes reduction、throughput-vs-latency 划分，记录在第 VII 部分；本页只点明每个调度阶段从哪里接入它。最相关的两个页面是 Cost Model Overview（三个 cost-class family 与每个 cycle number 背后的 Target clock wiring）和 Bundle-Aware Cost（供给 LHS comparator 的 MaxResourceCycles bundle cost 与 LatencyBetween dependency latency）。

注意 — Stage 3 的“cost”是 legality，而不是 latency。 与 Stage 1 不同，bundle packer 不最小化 cycle；它满足一个约束（slot fit）。它唯一的 latency 输入，是 BundlePacker::Feed 中用来计算最早 legal bundle 的逐 op RAW latency。latency hiding，也就是把长 edge 重叠在 compute 之下，已经由 Stage 1 在 HLO 层决定。重新实现者不应在 packer 中重新运行 latency-hiding heuristic；那会重复计算优化。

置信度摘要

主张	证据
Stage 1 LHS 由 `RunHloScheduler` 作为两个顺序 pipeline（base 然后 overlap）调用	`RunHloScheduler` @ `0x1096fac0`（Brkga / `final_scheduler` / `async_scheduling` strings）
LHS `RunImpl` 主体在所有变体间共享	单个编译出的 `RunImpl` @ `0x136321a0`
LHS comparator 是 22-key `ReadySetLt` chain，反向 drain	`FindAndExtractBestNodeAvailable` @ `0x13618880`；`ScheduleComputation` @ `0x1362eb60`
Stage 2 把 matmul bin-pack 进 `MxuSequence` 并预留 MRB	`AssignMxusForSequenceGroup` @ `0x10f753c0`；`MrbChainAllocator::ExtendMrbReservation` @ `0x10f58800`
Stage 2 commit 下游 Stage 3 读取的 latch index	`SetLatchIndices` @ `0x10f3b4c0`；`AllocateMrbEntriesAsFifo` @ `0x10f3ef80`
Stage 3 是前向 greedy earliest-legal-bundle list scheduling	`PackBundles` @ `0x10a30a20`；`GlobalBundlePacker::Pack` @ `0x10a86420`；`BundlePacker::Feed` @ `0x14021f20`
Stage 3b 是面向 hardware loop 的 modulo scheduler，带 RecMII/ResMII II search	`TPUScheduleDAGModulo::findSchedule` @ `0x13b1d7c0`；`calculateResourceMII` @ `0x13c0bee0`；`calculateLargestLatencyMII` @ `0x13c0b840`
Stage 1 通过 `MaxResourceCycles` / `LatencyBetween` 定价；通过 `TpuAsyncTracker` gate	`MaxResourceCycles` @ `0x1c89b9e0`；`LatencyBetween` @ `0x1c89f820`；`GetResourceHazardType` @ `0x110015e0`
逐代 slot legality 是四个 `TpuBundleRestrictions` 子类（Jellyfish/Pufferfish/Viperfish/Ghostlite）	仅存在四个 `*BundleRestrictions::SetLimits` 符号（`0x1c457a40` / `0x1c457d80` / `0x1c458360` / `0x1c458860`）；没有 `Trillium`/`dragonfish` restriction 符号
"ILP-LHS" flag 只替换 async classifier；ILP MIP 是独立 opt-in memory scheduler（不是 dead）	`EnableIlpLatencyHidingScheduler` gate 只重新分类 async op；`ILPMemoryScheduler::Run` @ `0x10acd020` 通过 `GetMemorySchedulerAlgorithm` @ `0x10abd6a0` case 6 到达（switch arms Default0..LocalOrder9，ILP=6，构造 `xla::ILPMemoryScheduler`）
HLO scheduling annotation 不流入 bundle packing	LLO packer 不读取 HLO annotation；二者操作互不相交的 IR

交叉引用

LatencyHidingScheduler Core — Stage 1 core：反向 drain loop、22-key comparator、预计算 async depth/height、memory-pressure retry loop。
LHS post_layout — 0.0.40 中实际 Stage 1 placement：two-pipeline RunHloScheduler 及其 config/estimator/tracker wiring。
LHS post_layout_pre_fusion — 同一 LHS 主体接到 dead "Pre main fusion" slot；由 config delta 记录。
LHS ILP variant — EnableIlpLatencyHidingScheduler async-classifier swap 和 dead ILPMemoryScheduler MIP。
ResourceType Taxonomy — gate Stage 1 overlap 的 0..46 ResourceType enum、hazard class，以及 AsyncTracker → core registry。
MxuSequence / SequenceInfo — Stage 2 每序列记录和 set_mxu commit。
MXU Assignment Bin-Packer — Stage 2 AssignMxusForSequenceGroup bin-packing 算法。
MRB Chain Allocator — Stage 2 accumulation-chain reservation timeline 与 jitter model。
MRB FIFO / MSR Placement — Stage 2 result-FIFO 与 matrix-shift-register placement。
Latch Assignment & Overrun — Stage 2 latch-index assignment 与逐代 overrun handshake。
LLO → Bundle Packing — Stage 3 前向 greedy packer、逐代 slot matrix、branch/barrier/phi handling。
Bundle Modulo Scheduling — Stage 3b 面向 hardware loop 的 II-search software pipelining。
Per-Gen Encoder Latch Serialization — Stage 2 的 latch field 如何序列化进逐代 bundle word。
Cost Model Overview — 每个 cycle number 背后的三个 cost-class family 和 Target clock wiring。
Bundle-Aware Cost — 供给 Stage 1 comparator 的 MaxResourceCycles bundle cost 和 LatencyBetween dependency latency。
Bundle Model Overview — Stage 3 打包、Stage 2 latch 序列化进入的 VLIW bundle word。
Binary: extracted/libtpu-0.0.40-cp314-cp314-manylinux_2_31_x86_64/libtpu/libtpu.so（build-id 89edbbe81c5b328a958fe628a9f2207d）
索引条目： Part VIII — Instruction Scheduling & Bundle Packing — 返回索引

TPU 调度 Pipeline ​

摘要 ​

阶段顺序 ​

变体映射 ​

消费内容 ​

Stage 2 — MXU Sequence Assignment 与 MRB Allocation ​

用途 ​

结构 ​

内部循环的 Modulo Scheduling（Stage 3b） ​

子页 ​

各阶段如何定价 ​

置信度摘要 ​

交叉引用 ​