LHS：ILP 变体

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

摘要

libtpu.so 带有一个字面上以整数线性规划调度命名的类 — xla::ILPMemoryScheduler — 以及名为 xla_tpu_enable_ilp_latency_hiding_scheduler 的标志族和一个六字段 protobuf：xla::jellyfish::IlpLatencyHidingSchedulerOptions。如果读者期待“LatencyHidingScheduler 的 ILP 变体”是一个由求解器支撑、替代贪心列表调度器的单一实现，就会被误导。静态分析显示，这个名称覆盖了两条结构上无关的代码路径；它们因共享上游 options 消息而命名相似，但接入的是两个不同的子系统。

第一条路径是按 computation 运行的内存调度 MIP — xla::ILPMemoryScheduler::Run（@0x10acd020）委托给 worker ILPMemorySchedulerForComputation::Run（@0x10acdf00）。这是一个真正的整数规划：它构建 operations_research::math_opt::Model，声明二元 placement/liveness 变量和一个连续 peak_memory 变量，生成十一组线性约束，最小化 peak_memory，并通过 OR-Tools math_opt::Solve 用 SolverType = CP_SAT (4) 求解。当 xla::GetMemorySchedulerAlgorithm（@0x10abd6a0）分派的 MemorySchedulerProto::Value 枚举被设为 ILP (6) 时，它就是被选中的内存调度器：case 6 构造该类（vptr off_217FA470），而该 dispatcher 会从实际 pipeline 启动路径调用（RunHloScheduler、PostMainFusionHloOptimize）。它是一个按需调度器，不是默认项 — DEFAULT (0) 选择 DefaultMemoryScheduler — 但它完全可达。

另一条路径驱动 latency-hiding scheduler，且不包含 ILP：在规范调度器启动 RunHloScheduler（@0x1096fac0）内部，xla::jellyfish::EnableIlpLatencyHidingScheduler（@0x1d6b7e00）在两个 async-op classifier lambda 之间选择。启用后，更宽的 lambda $_1 会向规范 xla::LatencyHidingScheduler 暴露更多 instruction kind 作为 async-overlap 候选 — 两种情况下运行的都是这个 pass。按 0.0.40 中的接线，“ILP-LHS”标志的含义是“暴露更多 async 候选”，而不是“切换到 ILP 调度器”。这个标志独立于选择内存调度 MIP 的 MemorySchedulerProto::Value::ILP 枚举：两者共享 IlpLatencyHidingSchedulerOptions 消息名，但不是同一个开关。

对于重新实现，契约是：

xla::ILPMemoryScheduler 对象布局（112 字节），以及其硬编码的 ≤ 8 HloValues 回退到 xla::BacktrackScheduler，同时受对象偏移 +104 处的 use_backtrack_fallback 标志控制。
完整 ILP 模型：三组二元变量（I_、Vs_、Ve_）加连续 peak_memory，十一组 AddLinearConstraint，单项 min peak_memory 目标，以及使用默认（无界）SolveParameters 的 CP-SAT 求解。
dispatcher 接线：GetMemorySchedulerAlgorithm 在 MemorySchedulerProto::Value switch 的 case ILP (6) 构造 ILPMemoryScheduler；只有当编译环境选择该枚举值时 MIP 才运行。
RunHloScheduler 内的 EnableIlpLatencyHidingScheduler gate，以及 $_1 vs $_2 classifier 语义 — 这是一个不触及 MIP 的独立开关。
IlpLatencyHidingSchedulerOptions proto：六个字段中哪一个被读取（一个），以及哪五个是惰性的。


MIP 外层 Run	`xla::ILPMemoryScheduler::Run(HloComputation*, ...)` @ `0x10acd020`
MIP worker Run	`(anon)::ILPMemorySchedulerForComputation::Run` @ `0x10acdf00`
MIP vtable	`0x217fa460`（vptr `off_217FA470`），typeinfo `0x217fa490` — 一个 consumer：dispatcher case 6
Dispatcher	`xla::GetMemorySchedulerAlgorithm` @ `0x10abd6a0` — case `ILP (6)` 构建 MIP
实际 gate	`xla::jellyfish::EnableIlpLatencyHidingScheduler` @ `0x1d6b7e00`
Gate 调用方	`(anon)::RunHloScheduler` @ `0x1096fac0`（gate 在 line 1084）
Async classifiers	`$_1`（ILP-on）@ `0x10977140`，`$_2`（regular）@ `0x10977420`
Options proto	`xla::jellyfish::IlpLatencyHidingSchedulerOptions`（6 个字段，1 个被读取）
求解器后端	仅 `math_opt::Solve` @ `0x10af9d20`，`SolverType = CP_SAT (4)`

一个名称背后的两条路径

标志族 xla_tpu_enable_ilp_latency_hiding_scheduler 并不选择调度器。追踪所有带有字符串 ILP 和 IlpLatencyHiding 的符号引用，会落到两个不相交的子系统，它们由两个独立开关到达：

text

two independent "ILP" switches
   │
   ├── (A) async-op classifier switch in canonical LHS
   │       flag xla_tpu_enable_ilp_latency_hiding_scheduler / proto +48
   │       EnableIlpLatencyHidingScheduler(env) @0x1d6b7e00
   │          └─ RunHloScheduler @0x1096fac0  →  picks lambda $_1 vs $_2
   │                └─ AddPass<xla::LatencyHidingScheduler>   (canonical, unchanged)
   │
   └── (B) per-computation memory-scheduling MIP
           MemorySchedulerProto::Value::ILP (6)   ← env field at +268
           GetMemorySchedulerAlgorithm @0x10abd6a0  case 6 → constructs ILPMemoryScheduler
              └─ xla::ILPMemoryScheduler::Run @0x10acd020
                    └─ ILPMemorySchedulerForComputation::Run @0x10acdf00
                          └─ math_opt::Solve(..., CP_SAT, ...) @0x10af9d20
```text

两条路径不共享状态、成本模型或入口点。路径 (B) 是唯一包含整数规划的路径；路径 (A) 扩大输入到贪心 LHS 的 async 表面。两者都可达，但通过*不同*旋钮：路径 (A) 通过 `enable_ilp_latency_hiding_scheduler` 标志/proto bool，路径 (B) 通过 env 偏移 +268 处的 `MemorySchedulerProto::Value` 枚举。把两者都命名为“ILP”是上游源码的构建期产物：切换路径 (A) 的 proto 字段位于同一个 options 消息中，而该消息也包含会调节路径 (B) 的（在此构建中惰性的）旋钮。

> **陷阱 —** `IlpLatencyHidingSchedulerOptions` proto 是一个单一消息，其*唯一*实际 consumer 是路径 (A) 的 classifier 开关（+48 处的 bool）。六个字段中的五个（`max_solver_deterministic_time`、`computation_size_threshold`、`use_ilp_schedule_sequence`、`also_minimize_total_lifetime`、`min_compute_latency`）描述路径 (B) 的求解器，但路径 (B) 一个也不读取 — 它从 `GetMemorySchedulerAlgorithm` 调用参数取得 `AliasInfo` 和 size function，而不是从这个 proto 取得。重新实现者如果把这个 options proto 接到 CP-SAT 求解器上，将无法复现二进制实际连接的行为：MIP 是由 `MemorySchedulerProto::Value` 枚举选择的，不是由这些字段选择的。

---

## 路径 (B)：`ILPMemoryScheduler` MIP

### 类层级与 dispatcher case

`xla::ILPMemoryScheduler` 是 `xla::ComputationSchedulerAlgorithm` 的子类，后者也是所有内存调度器共享的基类。它的 deleting destructor（`@0x10ad6900`）把基类 vptr 重置为 `ComputationSchedulerAlgorithm` vtable（`off_21CF7F08`）并释放 112 字节，确认了继承关系和对象大小。具体子类由 `xla::GetMemorySchedulerAlgorithm`（`@0x10abd6a0`）中的枚举分派选择，该分派对存储在 `env + 268` 处的 `MemorySchedulerProto::Value` 做 switch：

```c
// sub_10ABD6A0 — verbatim switch structure
switch ( *(_DWORD *)(a1 + 268) ) {
  case 0: ...  // DEFAULT       → xla::DefaultMemoryScheduler
  case 1: ...  // LIST          → xla::ListMemoryScheduler
  case 2: ...  // DFS           → xla::DFSMemoryScheduler
  case 3: ...  // POST_ORDER    → xla::PostOrderScheduler
  case 4: ...  // BRKGA         → xla::BrkgaMemoryScheduler    (vptr off_217FA2C0)
  case 5: ...  // BFS           → xla::BFScheduler
  case 6: ...  // ILP           → xla::ILPMemoryScheduler      (vptr off_217FA470)
  case 7: ...  // BACKTRACKING  → xla::BacktrackMemoryScheduler
  case 8: ...  // BRUTE_FORCE   → xla::BruteForceMemoryScheduler
  case 9: ...  // LOCAL_ORDER   → xla::LocalOrderScheduler (+ inner RandomOrderScheduler)
  default:
    LogMessage("hlo_scheduling_selector.cc", 62);
    CopyToEncodedBuffer("Unexpected memory scheduler: ", 29);
    return make_unique<xla::DefaultMemoryScheduler>(...);   // fallback
}

枚举名称和数值是从嵌入的 proto descriptor（MemorySchedulerProto.Value）按字节恢复的：DEFAULT=0, LIST=1, DFS=2, POST_ORDER=3, BRKGA=4, BFS=5, ILP=6, BACKTRACKING=7, BRUTE_FORCE=8, LOCAL_ORDER=9。Case 6 是 ILP，并构造 xla::ILPMemoryScheduler：它分配 112 字节，填充共享的 ComputationSchedulerAlgorithm 布局，然后把最终 vptr 设为 off_217FA470（ILP vtable 的第一个虚槽，vtable + 0x10），并把 +104 处的 bool 设为 1。默认分支在 hlo_scheduling_selector.cc:62 记录 "Unexpected memory scheduler: "，并回退到 DefaultMemoryScheduler。

注意 — 每个 case 1..8 都构建一个 112 字节对象，且具有完全相同的布局 — 基类 vptr off_21CF7F08、+8 处的 AliasInfo*、+16/+64 处的 size-function AnyInvocable、+88/+96 处的空 post-process std::function hook，以及 +104 处的 bool — 然后把最终 vptr patch 到子类 vtable。BRKGA（case 4）patch 到 off_217FA2C0；ILP（case 6）patch 到 off_217FA470。这些调度器是共享同一存储形状的同级类；只有 vptr 和 +104 bool 值不同。

对象布局与硬回退

外层 Run（@0x10acd020，约 861 行）具有如下重建签名：

cpp

StatusOr<HloInstructionSequence>
xla::ILPMemoryScheduler::Run(HloComputation* comp,
                             TuplePointsToAnalysis const& pts,
                             HloAliasAnalysis const& alias) const;
```text

| 偏移 | 字段（推断） |
|-------:|------------------|
| 0 | vptr → `ILPMemoryScheduler` vtable 第一个槽（`off_217FA470`，vtable base `0x217fa460`） |
| 8 | `AliasInfo const*` |
| 16 / 64 | size-function `AnyInvocable<int64_t(BufferValue const&)>`（RAII storage + live callable） |
| 48 | bool — 内联 `AnyInvocable` 的 alive marker |
| 88 / 96 | post-process `std::function<HloInstructionSequence(...)>` hook（默认空） |
| 104 | bool — `use_backtrack_fallback` |

回退测试在反编译中是逐字节精确的（外层 Run line 122）：

```c
// a3 = TuplePointsToAnalysis&;  a3+88 = values().size()
// a2 = this;  a2+104 = use_backtrack_fallback
if ( *(__int64 *)(a3 + 88) <= 8 && *(_BYTE *)(a2 + 104) )
    return xla::BacktrackScheduler::Run(this, &fallback_state, ...);   // greedy

只有当 （HloValues 数量 > 8）或（use_backtrack_fallback 被清除） 时才会走 ILP 路径。阈值 8 是硬编码的 — 它不是由 proto 字段 computation_size_threshold 驱动，后者存在但从未被读取。对于超过 8 个值的图，或回退开关被清除的图，MIP 无条件运行，没有 problem-size guard。

怪癖 — use_backtrack_fallback 成员（偏移 +104）是唯一控制回退的 bool，比较操作数 8 是指令流中的字面量。显然本应驱动这一点的 options-proto 字段（computation_size_threshold，int64）是惰性的。重新实现如果读取 computation_size_threshold 来决定回退规模，会偏离忽略它的二进制行为。

求解前模型构建（外层 Run）

// outer Run, line 168
operations_research::math_opt::Model::Model(&model, "ilp_memory_scheduling", 21);
// line 175 — the single continuous variable
ModelStorage::AddVariable(model_storage,
                          /*is_int=*/0, "peak_memory", /*name_len=*/11,
                          /*lb=*/0.0, /*ub=*/+inf);
// ... walk HloAliasAnalysis::buffers(), group HloValues by defining-instruction id
//     into absl::btree_map<HloValue const*, vector<HloInstruction*>, OrderHloValuesById>,
//     keeping only values whose defining instruction lives in *this* computation ...
// line 788 — delegate
ILPMemorySchedulerForComputation::Run(worker);
```text

模型名字符串 `"ilp_memory_scheduling"`（长度 21）和连续变量名 `"peak_memory"`（长度 11，`is_int = 0`）在反编译中都是字面量。每个 `HloValue` 的大小由偏移 64 处用户提供的 `AnyInvocable<int64_t(BufferValue const&)>` 定价 — 没有 `HighWaterMark` 风格的流式估算器；峰值是从 MIP 解本身恢复的。

---

## ILP 形式化（worker Run）

worker `ILPMemorySchedulerForComputation::Run`（`@0x10acdf00`，约 6764 行）构建并求解模型。下列所有锚点均已按字节对反编译验证；source-line number 指从 `LogMessage`/`AddSourceLocationImpl` 调用点恢复的 `platforms/xla/service/jellyfish/hlo_scheduling/memory_schedulers.cc`。

### 决策变量

三组变量全是二元变量，通过 `ModelStorage::AddVariable(model, /*is_int=*/1, name, name_len, /*lb=*/0, /*ub=*/1)` 构建。连续 `peak_memory`（外层 Run 中构建，`is_int = 0`）是第四个。

| 族 | 名称模式 | 反编译行 | 索引范围 | 含义 |
|--------|--------------|---------------|--------------|---------|
| `I_` | `"I_<inst>,<slot>"` | 817 / 843 | （instruction id，slot `0..N-1`） | 当且仅当 instruction `i` 放在 slot `t` 时，`I_i,t = 1` |
| `Vs_` | `"Vs_<value>,<slot>"` | 1685 / 1710 | （HloValue id，slot `0..N-1`） | 当且仅当 value `v` 到 slot `t` 为止*已开始 live* 时，`Vs_v,t = 1` |
| `Ve_` | `"Ve_<value>,<slot>"` | 1807 / 1833 | （HloValue id，slot `0..N-1`） | 当且仅当 value `v` 到 slot `t` 为止*已结束*时，`Ve_v,t = 1` |
| — | `peak_memory` | 外层 Run 175 | 标量 | 连续变量；由每个 slot 的 live total 给出下界 |

前缀字符串字面量 `"I_"`、`"Vs_"`、`"Ve_"` 由 `absl::StrCat` 发出，索引和 slot 使用 `absl::numbers_internal::FastIntToBuffer`。变量到 `Variable` 的映射是 instruction 使用的 `absl::flat_hash_map<std::pair<HloInstruction*,int>, math_opt::Variable>`，以及 value 使用的 `flat_hash_map<std::pair<long,int>, math_opt::Variable>`。`Vs_`/`Ve_` 将 liveness 建模为单调阶跃函数（累积的“started by”/“ended by”指示器），这使得每个 slot 的内存和可以写成线性表达式。

### 约束

worker 中恰好有**十一个** `Model::AddLinearConstraint` 调用点，位于下列按字节验证的行。功能分组是从周围循环结构和各站点迭代的变量族推断的；逻辑陈述是重建结果，不是字面系数恢复。

| # | 行 | 族 | 逻辑陈述 |
|---|------|--------|-------------------|
| 1 | 1362 | per-instruction | `Σ_t I_i,t == 1` — 每个 instruction 恰好位于一个 slot |
| 2 | 1573 | per-slot | `Σ_i I_i,t == 1` — 每个 slot 恰好容纳一个 instruction |
| 3 | 2252 | start init | `Vs_v,0 == I_def(v),0` — start CDF 由定义 instruction 初始化 |
| 4 | 2584 | start step | `Vs_v,t − Vs_v,t-1 == I_def(v),t` — monotone start 在定义 slot 阶跃 |
| 5 | 2952 | start coupling | `v` 的 user 不能在 `v` starts 之前运行 |
| 6 | 3356 | end step | `Ve_v,t − Ve_v,t-1 ≥ I_last_use(v),t` — end CDF 在 last-use slot 阶跃 |
| 7 | 3583 | end ≤ start | `Ve_v,t ≤ Vs_v,t` — value 不能在 start 前结束 |
| 8 | 4084 | precedence (data) | 对每条 HLO data edge，`Σ_{t'≤t} I_pred,t' ≥ I_succ,(t+1)` |
| 9 | 4483 | precedence (control) | control dependencies 使用相同形状 |
| 10 | 4887 | per-slot memory | 对每个 slot `t`，`Σ_v size(v)·(Vs_v,t − Ve_v,t) ≤ peak_memory` |
| 11 | 5491 | objective coupling | 将 `peak_memory` 与有界 per-slot expression 闭合 |

站点 1–2（`I_` placement 约束）是 assignment-problem 核心 — instruction 到 slot 的一个排列。站点 3–5 迭代 `Vs_` map，站点 6–7 迭代 `Ve_` map（两者都是双层循环：外层 HloValue，内层 slot）。站点 8–9 迭代按 `OrderHloValuesById` 分组的 predecessor/user map。站点 10 是 liveness 到 memory 的归约。

> **注意 —** 每个 `LinearExpression` 构造器内部的精确系数向量和 RHS 值尚未恢复：构造器被内联，并带有匿名 `flat_hash_map<Variable, double>` 参数，需要逐调用点符号检查。变量*族*、*数量*（十一个）和调用点*行号*是确定的；上面对每个约束的精确代数是中等置信度的重建。

### 目标函数

目标是单个连续变量，取最小值（worker Run lines 5567–5573）：

```c
v1196 = 0x3FF0000000000000LL;        // IEEE-754 double 1.0  — peak_memory coefficient
LinearExpression::LinearExpression(expr, /*terms=*/{peak_memory}, /*n=*/1, /*offset=*/0.0);
Model::SetObjective(model, expr, /*maximize=*/0);   // 0 == minimize

min peak_memory，系数 1.0，offset 0.0，maximize = false。没有第二目标：proto 字段 also_minimize_total_lifetime 从未被读取。

日志

worker 在 memory_schedulers.cc:283 发出一次性模型摘要：

text

Computation: <name>
Number of instructions: <N>
Number of variables: <model.Variables().size()>
Number of constraints: <model.LinearConstraints().size()>
Number of values related to computation: <values.size()>
```text

字符串字面量（`"Computation: "` len 13，`"Number of variables: "` len 21，`"Number of constraints: "` len 23，`"Number of values related to computation: "` len 41）已在 worker lines 5594/5606/5611/5616 按字节验证。成功求解时，它在 line 300 记录 `"Peak memory: "`；在 `VLOG ≥ 2` 时，它转储每个 instruction 的 `" scheduled at "`（line 423）以及每个 value 的 `Value defined at ... with size : ...`（line 454）。

---

## 求解器后端、参数、终止

### CP-SAT 是唯一后端

`Solve` 通过 `math_opt::Solver::New`（`@0x10b259a0`）分派，后者验证 init args 和 model，然后调用 `AllSolversRegistry::Create(registry, solver_type, model, init_args)`。registry 只填充一次，发生在 `_GLOBAL__sub_I_cp_sat_solver.cc`（`@0x212ca300`）的静态初始化中：

```c
v4[0] = operations_research::math_opt::CpSatSolver::New;
operations_research::math_opt::AllSolversRegistry::Register(registry, /*solver_type=*/4u, v4);

AllSolversRegistry::Register（@0x10b273c0）只有这一个调用方。SolverType = 4 是 SOLVER_TYPE_CP_SAT。任何其他 SolverTypeProto 值都会解析为缺失的 registry entry，并返回失败 status。MIP 的 solve site 字面传入 4（worker line 5833）：

operations_research::math_opt::Solve(result, model, /*solver_type=*/4, args, init_args);
```text

### 参数：没有应用边界

求解前构造的 `SolveArguments` 块（worker lines ~5780–5832）是零初始化的：空 message-callback `std::function`、空 solver-callback `std::function`、null `SolveInterrupter`。`SolveParametersProto` 保持默认值 — `enable_output = false`（只在 `VLOG` 下提升）、时间无界、没有 iteration/node limit。各 per-solver 参数结构（`SatParameters`、`GScipParameters`、`GlopParameters` 等）会被构造，但只有 CP-SAT 的 `SatParameters` 被消费，且它保持默认。

> **陷阱 —** 标志 `xla_tpu_ilp_scheduler_max_solver_deterministic_time`（data `@0x223c1db0`，注册于 `tpu_compilation_environment.cc`）和 proto 字段 `max_solver_deterministic_time` 都有存储值，但 worker Run 中**两者都不读取**。`SatParameters` deterministic-time 和 wall-time limit 保持默认值（CP-SAT 默认：无限制）。求解器会运行到完成或硬失败 — 没有收敛边界。

### 终止

```c
SolveResult r = Solve(...);
if ( !r.ok() )                                  // → memory_schedulers.cc:297
    return r.status().AddSourceLocation(...);
if ( !r->termination.EnsureIsOptimalOrFeasible().ok() )   // → memory_schedulers.cc:298
    return ...;
// success: read peak_memory and every I_i,t via SolveResult::variable_values()

EnsureIsOptimalOrFeasible（worker line 5872）接受 OPTIMAL 和 FEASIBLE — 即使不是最优，也会消费任何可行 primal。INFEASIBLE、UNBOUNDED、IMPRECISE 和 NO_SOLUTION_FOUND 都会在 source line 298 转换为失败的 absl::Status。随后调度由 RetrieveSchedule 恢复：对每个记录的 (instruction, slot) pair，读取变量的 optimal value，构建 std::pair<long /*slot*/, HloInstruction*>，并按 slot 用 std::__stable_sort（实例化 @0x10adf140 / @0x10adf3a0 / @0x10adf5e0）排序，以发出最终的 HloInstructionSequence。

路径 (A)：实际的 async-op classifier 开关

EnableIlpLatencyHidingScheduler（@0x1d6b7e00）的唯一调用方是 RunHloScheduler（@0x1096fac0）。gate 位于反编译 line 1084：

if ( EnableIlpLatencyHidingScheduler(env) ) {
    is_async_pred = &RunTensorCoreAsyncOpScheduler::$_1;   // 0x10977140 — wider
} else {
    is_async_pred = &RunTensorCoreAsyncOpScheduler::$_2;   // 0x10977420 — narrower
    use_full_kind_set = EnableSchedulingAnnotationPropagation(env);   // line 1113
}
// both paths land here:
pipeline.AddPass<xla::LegalizeSchedulingAnnotations>(cfg);   // line 1136
pipeline.AddPass<xla::LatencyHidingScheduler>(ctx, core);    // line 1137 — canonical
pipeline.AddPass<xla::ConstantDeferring>();                  // line 1138
```text

规范 `xla::LatencyHidingScheduler` 在**两种**情况下都会运行（line 1137）。gate 位于 `AddPass` 之前，因此 scheduler pass 本身无法知道使用了哪个 classifier — 只有它接收的 async-candidate set 发生变化。参见 [LatencyHidingScheduler 核心](latency-hiding-scheduler-core.md)。

### Classifier `$_1`（ILP-on，`@0x10977140`）

从反编译按字节验证：

```c
// is_async($_1)(inst):
attr = inst->get_frontend_attribute("keep_original_sequence_order_in_group", 37);
if ( attr.present && attr.value_len == 4 && *(uint32*)attr.value == 1702195828 )  // "true"
    return true;                       // honour user-pinned ordering
return inst->IsOutputFusion()
    || inst->IsLoopFusion()
    || inst->opcode() == 0x82
    || inst->IsCustomCall("tpu_custom_call")
    || inst->IsCustomFusion()
    || inst->opcode() == 40;           // kDot

字面量 1702195828 是 ASCII 字节 "true" 的 little-endian dword；value-length guard（4）是与 4 字节字符串区域的比较。37 字符的 attribute name 已 intern。

Classifier `$_2`（regular，`@0x10977420`）

// is_async($_2)(inst):
return inst->IsOutputFusion()
    || inst->IsCustomCall("tpu_custom_call")
    || inst->IsCustomFusion()
    || inst->opcode() == 40;           // kDot
```text

差异正好是：`$_1` 额外把 `IsLoopFusion` 和 opcode `0x82` 视为 async，并遵守 `keep_original_sequence_order_in_group="true"` pin。因此 ILP-LHS 标志会扩大规范调度器看到的 async-overlap 表面 — 这会改变 `DefaultSchedulerCore` 使用的 recurrence latencies — 但不会改变调度算法。成本模型是常规 LHS 成本模型（`xla::DefaultSchedulerCore` + `xla::LatencyEstimator`）；参见 [调度器概览](overview.md) 和 [成本模型概览](../cost/overview.md)。

---

## `IlpLatencyHidingSchedulerOptions` Proto

该消息已注册，可从标志解析，并可通过 `GetIlpLatencyHidingSchedulerOptions`（`@0x1d6b7e60`）到达。它的 `Clear`（`@0x1db24ea0`）会清零 struct 偏移 +16 处的 6-bit presence bitmap（`& 0x3F`）以及覆盖偏移 `0x18..0x32` 的 payload 区域；`enable_ilp_latency_hiding_scheduler` bool 位于偏移 +48（`0x30`），处在该 payload 内。

```protobuf
message IlpLatencyHidingSchedulerOptions {
  optional bool   enable_ilp_latency_hiding_scheduler = 1;  // C++ struct +48 — LIVE
  optional double max_solver_deterministic_time      = 2;   // inert
  optional int64  computation_size_threshold         = 3;   // inert
  optional bool   use_ilp_schedule_sequence          = 4;   // inert
  optional bool   also_minimize_total_lifetime       = 5;   // inert
  optional double min_compute_latency                = 6;   // inert
}

实际 consumer 只读取 +48 处的 bool（gate @0x1d6b7e00）：

bool EnableIlpLatencyHidingScheduler(env) {
    auto opts = GetIlpLatencyHidingSchedulerOptions(env);
    char field48 = ((char*)&opts)[48];               // enable_ilp_latency_hiding_scheduler
    ~opts.~IlpLatencyHidingSchedulerOptions();
    if ( field48 ) return true;                       // proto override on
    AutoOr<bool> f = AutoOr<bool>::FromProtoOrDie(env + 1608);  // else fall back to the flag
    return (~f & 0x101) == 0;
}
```text

proto override 或 `xla_tpu_enable_ilp_latency_hiding_scheduler` absl flag（data `@0x223c1d50`）任一都可把 classifier 切换到 `$_1`。`(~v & 0x101) == 0` idiom 是 `AutoOr<bool>` unwrap（value bit 和 present bit 都置位）。

| 字段 | 在 0.0.40 中读取？ | 效果 |
|-------|-----------------|--------|
| `enable_ilp_latency_hiding_scheduler` (+48) | 是 | 切换规范 LHS 的 async classifier：`$_2 → $_1` |
| `max_solver_deterministic_time` | 否 | 无 — `SatParameters` 保持默认 |
| `computation_size_threshold` | 否 | 无 — 回退阈值硬编码为 `8` |
| `use_ilp_schedule_sequence` | 否 | 无 |
| `also_minimize_total_lifetime` | 否 | 无 — 单项目标 |
| `min_compute_latency` | 否 | 无 |

五个惰性字段只出现在生成的 reflection code 中（`_table_` `@0x21cfa308`、`Clear`、`MergeImpl`、`_InternalSerialize`、`ByteSizeLong`、`CopyFrom`、`InternalSwap`、`GetMetadata`、`GetClassData`）。任何非生成函数都没有调用它们的 `set_*`/`get_*` accessor。

### 标志

| 标志（data symbol） | 0.0.40 中可达的 consumer |
|--------------------|------------------------------|
| `FLAGS_xla_tpu_enable_ilp_latency_hiding_scheduler` (`@0x223c1d50`) | `EnableIlpLatencyHidingScheduler`（通过 `AutoOr<bool>`） |
| `FLAGS_xla_tpu_ilp_latency_hiding_scheduler_options` (`@0x223c4750`) | `GetIlpLatencyHidingSchedulerOptions`（通过 `AutoOr<message>`） |
| `FLAGS_xla_tpu_ilp_scheduler_max_solver_deterministic_time` (`@0x223c1db0`) | **无** — 已注册但未读取 |

---

## 相对于贪心列表调度器的替换条件

给重新实现者的直接答案是：在 `libtpu-0.0.40` 中，MIP **只有在编译环境设置 `MemorySchedulerProto::Value::ILP (6)` 时**才替代贪心内存调度器 — 它是 opt-in，不是默认值。

- MIP 路径 (B) 在 `GetMemorySchedulerAlgorithm` 分派到 case 6 时替代贪心内存调度器。默认值是 `DEFAULT (0)` → `DefaultMemoryScheduler`，因此未配置的构建永远不会运行 MIP。贪心 `BacktrackScheduler` 仍可从路径 (B) 到达，但只作为 ILP 调度器自身 `Run` 内的 `≤ 8 HloValues` 回退。
- 实际路径 (A) — `enable_ilp_latency_hiding_scheduler` 标志 — **完全不会**替代列表调度器。它保留规范 `xla::LatencyHidingScheduler` pass，只扩大输入给它的 async-candidate set。贪心 `DefaultSchedulerCore` walk 不变；只有输入变化。路径 (A) 与路径 (B) 的枚举选择正交。

MIP 是一个完整、可达的调度器 — `Model`、十一条约束、`min peak_memory`、CP-SAT 求解 — 由一个枚举值选择。它不是默认值，也未在默认 pipeline 运行中观察到，但它已接入 dispatcher，并由 case 6 构造。

> **注意 —** `xla::ILPMemoryScheduler` 是 `libtpu.so` 中唯一的内存调度 MIP，也是 OR-Tools `math_opt::Solve` consumer 之一（另外还有 auto-sharding 的 `FormulateAndSolveMIPFromProblem` 和 MSA ILP pass）。在内存调度器中，它是唯一由求解器支撑的选项；贪心 `DefaultMemoryScheduler` 是默认项，也是大多数编译采用的路径。

---

## 函数与符号映射

| 符号 | 地址 | 角色 |
|--------|---------|------|
| `xla::ILPMemoryScheduler::Run` | `0x10acd020` | 外层 Run：布局、`≤8` 回退、模型 bootstrap |
| `(anon)::ILPMemorySchedulerForComputation::Run` | `0x10acdf00` | worker：变量、11 条约束、目标、求解 |
| `xla::ILPMemoryScheduler::~ILPMemoryScheduler` | `0x10ad6900` | deleting dtor；重置 base vptr |
| vtable `xla::ILPMemoryScheduler` | `0x217fa460` | vptr `off_217FA470`，typeinfo `0x217fa490`；一个 consumer（dispatcher case 6） |
| `xla::GetMemorySchedulerAlgorithm` | `0x10abd6a0` | enum dispatch；case `ILP (6)` 构建 MIP |
| `xla::jellyfish::EnableIlpLatencyHidingScheduler` | `0x1d6b7e00` | 实际 gate；读取 proto +48，然后读取 flag |
| `xla::jellyfish::GetIlpLatencyHidingSchedulerOptions` | `0x1d6b7e60` | options accessor |
| `(anon)::RunHloScheduler` | `0x1096fac0` | gate 调用方；gate 位于 line 1084 |
| classifier `$_1`（ILP-on） | `0x10977140` | 更宽的 async predicate + sequence-order pin |
| classifier `$_2`（regular） | `0x10977420` | 更窄的 async predicate |
| `IlpLatencyHidingSchedulerOptions::Clear` | `0x1db24ea0` | 6-bit presence mask，payload `0x18..0x32` |
| `math_opt::Solve` | `0x10af9d20` | CP-SAT 求解入口 |
| `math_opt::Solver::New` | `0x10b259a0` | 基于 `SolverType` 的 registry dispatch |
| `math_opt::AllSolversRegistry::Register` | `0x10b273c0` | 唯一调用方只注册 CP-SAT |
| `math_opt::CpSatSolver::New` | `0x10adff20` | 唯一注册后端 |
| `_GLOBAL__sub_I_cp_sat_solver.cc` | `0x212ca300` | static-init：`Register(..., 4u, CpSatSolver::New)` |
| `RetrieveSchedule` stable-sort | `0x10adf140` / `0x10adf3a0` / `0x10adf5e0` | 以 slot 为 key 合并求解出的 schedule |

---

## 交叉引用

- [LatencyHidingScheduler 核心](latency-hiding-scheduler-core.md) — 路径 (A) 输入的规范 scheduler pass；唯一运行的 LHS 算法
- [LHS：post_layout / final 变体](lhs-post-layout.md) — 同一 pipeline 中最终的 post-layout scheduling 变体
- [调度器概览](overview.md) — scheduling 位于 lowering 与 encoding 之间的位置
- [ResourceType 分类](scheduler-resourcetype-model.md) — 调度器的 47-ID 并发模型，不同于此 MIP 的 peak-memory 目标
- [成本模型概览](../cost/overview.md) — 路径 (A) 的规范 scheduler 使用的 `GetLatencyBetween` 背后的 HLO-level 成本模型（此 MIP 只使用 peak-memory size function，不使用成本模型）
- [返回索引](../index.md) — Part VIII — Instruction Scheduling & Bundle Packing

LHS：ILP 变体 ​

摘要 ​

一个名称背后的两条路径 ​

对象布局与硬回退 ​

求解前模型构建（外层 Run） ​

日志 ​

路径 (A)：实际的 async-op classifier 开关 ​

Classifier $_2（regular，@0x10977420） ​