LHS post_layout_pre_fusion 变体
本页中的所有地址都适用于来自
libtpu-0.0.40-cp314-cp314-manylinux_2_31_x86_64wheel 的libtpu.so（build-id89edbbe81c5b328a958fe628a9f2207d）。由未剥离、完整符号 ELF 的静态分析重建（nm -C可解析每个方法）；其他版本会有所不同。
摘要
post_layout_pre_fusion 变体是 LatencyHidingScheduler (LHS) 在 HloOptimizeAfterLayoutAssignment (0x10962660) 的 "Pre main fusion" 槽位中本应呈现的形态，即在布局分配之后、但在专用 "Main fusion" 流水线之前运行的调度器，从而调度仍未融合的 op，并让产生的重叠画像可以反馈给融合 pass。本页的决定性发现是一个否定结论，对任何重新实现者都很关键：在 v0.0.40 中，这个槽位没有接线。 HloOptimizeAfterLayoutAssignment 第 1169 行构建的 "Pre main fusion" HloPassPipeline 不包含 AddPass<LatencyHidingScheduler>，且 HloOptimizeAfterLayoutAssignment 从不调用 RunHloScheduler。所有可达的 AddPass<LatencyHidingScheduler> 都位于 RunHloScheduler (0x1096fac0) 中，这是第 7 阶段，即主融合之后，也是规范的 post_layout pass。
该槽位失效的原因是结构性的，而非偶然。共享的 RunImpl (0x136321a0) 一开始就在 !module->has_schedule() 上触发 LogMessageFatal："LatencyHidingScheduler expects a base schedule that minimizes memory pressure." 这个基础 schedule 由 HloMemorySchedulerWithBrkgaFallback 产生，而它只在 RunHloScheduler 内部添加（第 567 行）。在融合前槽位没有先前的内存 schedule，因此那里若存在可运行的 LHS，就会中止编译。因此 "Pre main fusion" 流水线不能承载规范的 OSS LHS；它实际运行的、唯一与 collective/schedule 相邻的工作是 CollectivePipeliner 及其清理步骤，其中没有任何一个添加 LHS。
本页以 wiki 记录另外两个变体的方式记录此变体：说明它与可运行 post_layout pass 以及共享 scheduler core 的差异。由于 pass 主体 (RunImpl)、配置构建器 (GetSchedulerConfig)、估算器选择器 (GetLatencyEstimator) 和 tracker 构建器 (GetTpuAsyncTracker) 都是与 post_layout 共享的单个已编译实例，差异完全在于融合前调用点会传入的输入，而在此构建中没有调用点提供这些输入。重新实现者若移植未来某个确实接线该槽位的 libtpu，只需要三个差异：融合前的 memory_limit_bytes、清零的 v45 post-SC-assignment 重跑字节，以及空的 reserved-SparseCore 集合。
对于重新实现，约定如下：
- 该槽位在 v0.0.40 中是死的。
"Pre main fusion"流水线存在（0x10962660第 1169 行），但未添加 LHS；仅有的两个AddPass<LatencyHidingScheduler>位置（RunHloScheduler第 1137 / 1411 行）都在第 7 阶段。对整个反编译结果 grep 规范的AddPass<LatencyHidingScheduler>模板，只会在RunHloScheduler内找到引用。 - 原因是
has_schedule()FATAL。RunImpl第 280 行是LogMessageFatal；没有先前的HloMemorySchedulerWithBrkgaFallback基础 schedule（在RunHloScheduler第 567 行、第 7 阶段添加），LHS 会中止而不是 no-op。融合前槽位没有这样的基础 schedule。 - pass 主体是共享的
RunImpl(0x136321a0)，与post_layout完全相同：has_schedule()前置条件、computation 列表过滤、调度前/后的LatencyHidingStatisticsVLOG(1) pass，以及带0.9松弛因子的、感知碎片的内存压力重试循环。 - 相对
post_layout的配置差异是三个调用点值，而不是不同的构建器：v45 == 0（没有 SC-assignment 重跑）、融合前memory_limit_bytes，以及空的 reserved-SparseCore 集合。137 字节SchedulerConfigschema、估算器分支和资源模型在其他方面都与post_layout逐字节相同。
| 预期阶段 | 第 6 阶段 — HloOptimizeAfterLayoutAssignment 的 "Pre main fusion" 子流水线（0x10962660，第 1169 行的字符串） |
| 在 v0.0.40 中已接线？ | 否 — 融合前流水线中没有 AddPass<LatencyHidingScheduler>；该槽位是死的 |
| 为何失效 | RunImpl has_schedule() LogMessageFatal（0x136321a0 第 280 行）；融合前不存在基础 schedule |
| 实际可运行的 LHS | RunHloScheduler (0x1096fac0) 中的 post_layout，第 7 阶段，主融合之后 |
| 共享 pass 主体 | xla::LatencyHidingScheduler::RunImpl — 0x136321a0（1217 行反编译；与 post_layout 相同） |
| 共享配置构建器 | xla::jellyfish::(anon)::GetSchedulerConfig — 0x10974aa0（137 行反编译，137 字节 POD） |
| 相对 post_layout 的配置差异 | v45 == 0；融合前 memory_limit_bytes；空的 reserved-SparseCore 集合 |
| 规范 AddPass 模板 | 0x10975c40 — 只从 RunHloScheduler 引用 |
| 置信度 | 高（否定结论在两个函数中都有字节锚点） |
槽位位置
目的
HloOptimizeAfterLayoutAssignment (0x10962660，2651 行反编译) 是第 6 阶段驱动器。它在 LayoutAssignment（第 5 阶段）和最终调度（第 7 阶段）之间构建一系列具名 HloPassPipeline 对象。融合前 LHS 本应添加到其 "Pre main fusion" 流水线，也就是名称字符串存储在第 1169 行的那条流水线。
阶段布局（恢复）
text
DeepseaCompilerBase::RunHloPasses (0x1093a420)
...
RunHloPassesAfterLayoutAssignment (0x21395d20)
└─ HloOptimizeAfterLayoutAssignment (0x10962660) [Phase 6]
"Post layout assignment" HloPassPipeline (SSO-built; tail "signment" at line 707)
"Pre main fusion" HloPassPipeline (line 1169) ← THIS SLOT (no LHS added)
CreateNestedHloPipeline "layout fixup" (line 1176)
CreateNestedHloPipeline "host offloading" (line 1396)
CollectivePipeliner + PipeliningCleanups (collective-adjacent; no AddPass<LHS>)
(60+ inline AddPass<...> sites — none is LatencyHidingScheduler)
"Main fusion" HloPassPipeline (line 2279)
PostMainFusionHloOptimize (0x10966560) [Phase 6 tail; called at line 2614]
PostOptimizationPipeline (0x1093fd40) [Phase 7 driver; called from RunHloPasses line 1987]
└─ RunHloScheduler (0x1096fac0) [Phase 7] ← the live LHS
```text
> **注 —** `"Post layout assignment"` 流水线名称并不是单个 `.rodata` C 字符串；它由内联 16 字节 SSO 存储组装而成，其尾部 `"signment"` 是 `0x10962660` 第 707 行的字面量（`strcpy(&v342[14], "signment")`）。同样的栈构建字符串手法也隐藏了 `RunHloScheduler` 中的流水线名称（`final_scheduler` / `async_scheduling`）；见 [`post_layout`](lhs-post-layout.md)。重新实现者若在二进制中 grep `"Post layout assignment"`，找不到它。
### 字节锚定的否定结论
```c
// HloOptimizeAfterLayoutAssignment @ 0x10962660
// "Pre main fusion" pipeline built at line 1169.
// rg over this function for LatencyHidingScheduler / AddPass<...Latency...>
// / RunHloScheduler → NO MATCHES.
// It builds the pipeline names and 60+ pass adds but never the LHS,
// and never invokes RunHloScheduler (Phase 7 is a sibling call, not nested here).RunHloScheduler 并不是从 HloOptimizeAfterLayoutAssignment 内部到达的；它是一个独立的第 7 阶段调用，由 PostOptimizationPipeline (0x1093fd40) 驱动，而 RunHloPasses 在第 1987 行调用它，也就是第 6 阶段的 HloOptimizeAfterLayoutAssignment 调用返回之后。规范的 AddPass<LatencyHidingScheduler> 模板主体 (0x10975c40) 在整个反编译中只从一个函数引用，即 RunHloScheduler，位于两个互斥位置（第 1137 和 1411 行），两者都在 "Main fusion" 之后。不存在可配置的融合前实例化。
槽位为何失效 — has_schedule() FATAL
目的
可运行 LHS 不能位于融合前的结构性原因，是 RunImpl 在日志之后做的第一件事：断言基础 schedule 存在。这是硬性的 LogMessageFatal，不是可恢复状态。
算法
c
// xla::LatencyHidingScheduler::RunImpl(module, exec_threads) // 0x136321a0
// line 280 / source latency_hiding_scheduler.cc:4110
if (!module->has_schedule()) {
LogMessageFatal(".../latency_hiding_scheduler.cc", 4110, "module->has_schedule()")
<< "LatencyHidingScheduler expects a base schedule that minimizes memory pressure.";
// ^ aborts the compile
}
```text
该前置条件要求的基础 schedule 由 `HloMemorySchedulerWithBrkgaFallback` 构建，它在 `RunHloScheduler` 第 567 行构造，位于第 7 阶段的 `final_scheduler` 流水线内部。`"Pre main fusion"` 流水线中没有等价步骤。因此在融合前添加 LHS 会在第一个 computation 上 FATAL。这与可运行 [`post_layout`](lhs-post-layout.md) pass 文档中的前置条件相同；在这里，它正是该变体*只是*一种配置、从未成为可运行 pass 的原因。
> **陷阱 —** FATAL 正是为什么“调度未融合 op 以反馈融合”不能通过简单地把 LHS 重新添加到更早流水线来实现。未来的 libtpu 若想要融合前 LHS，必须*同时*在同一个融合前流水线中在它之前添加内存调度器（BRKGA-fallback 基础调度器或等价物）；否则 `has_schedule()` CHECK 会中止。重新实现者不能只移动 `AddPass` 就满足该意图。
---
## Pass 主体 — 共享 `RunImpl`
### 目的
当（在未来构建中）某个融合前调用点*确实*添加 LHS 时，它调用的是与 `post_layout` 使用的同一个已编译 `RunImpl` (`0x136321a0`)。二进制中只有一个 `RunImpl`，即未修改的上游 `latency_hiding_scheduler.cc` 主体。该变体改变的是传给 `AddPass` 的配置和上下文，而不是主体。完整算法（drain loop、22 键比较器、异步资源模型）记录在 [scheduler core](latency-hiding-scheduler-core.md) 页面；本节钉住包围它的 `RunImpl` 外壳。
### 算法 — RunImpl 外壳
```c
absl::StatusOr<bool> RunImpl(HloModule *module,
const flat_hash_set<string_view>& exec_threads) { // 0x136321a0
VLOG(5) << "Original module:" << module->ToString(); // line ~270
// (1) precondition — FATAL if no base schedule (the dead-slot reason above)
CHECK(module->has_schedule()); // line 280, src 4110
// (2) computations to schedule: post-order, keep top-level OR force_delay group
to_schedule = MakeComputationPostOrder(exec_threads)
.filter(c -> c->caller_instructions().empty()
|| c->frontend_attribute("scheduling_group") == "force_delay");
VLOG(2) << "Computations to schedule " << to_schedule.size() // line 578
<< " size HloGraphNode: " << ... << " size HloEdge " << ...;
if (to_schedule.empty()) return false;
// (3) core init (+ real-cost-model init if config.use_real_cost_model)
scheduler_core_->SchedulerInitialize(module);
// (4) pre-schedule statistics (VLOG(1) only)
if (VLOG_IS_ON(1))
for (c : to_schedule)
LOG(INFO) << "] Statistics before scheduling:" // line 673
<< LatencyHidingStatistics(c, ...).ToString();
// (5) memory-pressure retry loop (see below)
// (6) final-limit log + (7) post-schedule statistics (VLOG(1))
LOG(INFO) << " LatencyHidingScheduler current memory usage: " // line 1088
<< (frag + computed_use()) << " bytes. Current limit: " << memory_limit();
if (VLOG_IS_ON(1))
for (c : to_schedule)
LOG(INFO) << "] Statistics after scheduling:" // line 1146
<< LatencyHidingStatistics(c, ...).ToString();
// (8) optional ScheduleProto dump
return true; // always reports "changed"
}调度前/后的 LatencyHidingStatistics pass（第 673 / 1146 行）使 LHS 可观测：在它们报告的 schedule 之前发出时，呈现的是输入基础 schedule 的浪费周期；之后的差值就是被隐藏的延迟。两者都会构建新的 BufferInfoTracker + ModulePressureState。它们受 VLOG(1) 控制，因此生产环境中除非 --v=1，否则没有成本。
注 —
force_delay过滤器（步骤 2）只在调用者带有scheduling_groupfrontend attribute"force_delay"时，才会把非顶层 computation 保留在 schedule 集合中。这与 core 的逐 computation 入口 RE2 匹配的字符串相同；它是让 callee computation 独立调度而不是与 caller 内联调度的唯一机制。
内存压力重试循环
c
// RunImpl memory retry @ 0x136321a0 (lines ~840-1054)
int64_t frag = config.estimate_fragmentation_size
? EstimateFragmentationSize(module, alias_analysis, alias_info) // line 840
: 0;
int64_t mem_limit = scheduler_core_->memory_limit();
for (int attempt = 0; attempt < scheduler_core_->memory_limit_retries(); ++attempt) {
if (frag + scheduler_core_->computed_use() <= mem_limit) break;
LOG(INFO) << "LatencyHidingScheduler current memory usage: " // line 874
<< (frag + computed_use())
<< " bytes, does not fit in initial limit: " << mem_limit // line 878
<< ". Setting the new limit to " << next_memory_limit(); // line 881
scheduler_core_->SchedulerInitialize(module);
mem_limit = (int64_t)scheduler_core_->next_memory_limit(); // relax-by-factor
scheduler_core_->set_memory_limit(mem_limit);
reschedule_all_computations();
if (config.estimate_fragmentation_size)
frag = EstimateFragmentationSize(...); // line 1054
}
```text
`next_memory_limit()` 将当前限制乘以 AVX 常量 `qword_A2DFD10`（第 892 / 918 行读取：`vmovsd xmm1, cs:qword_A2DFD10` / `vmulsd`），core 页面将其解析为 **`0.9`**，即每次重试将预算*缩小* 10%，把比较器的内存压力键（6-8）推到重叠键之上。在**融合前配置中，这个循环会是惰性的**，原因与第一次 `post_layout` 运行时相同：除非 `EnableSchedulerMemoryPressureTracking` 为 true，否则 `memory_limit == -1`。融合前配置永远不能设置 `v45` 重跑字节（按构造，它不是 post-SC-assignment 重跑），因此该循环在融合前触发的*唯一*方式是压力跟踪标志。
> **陷阱（低）—** 精确的 `qword_A2DFD10 = 0.9` 十进制值是在 [scheduler core](latency-hiding-scheduler-core.md) 页面从 IEEE 位模式解析得到的；`EstimateFragmentationSize`（TPU 私有；是一个 `GlobalDecreasingSizeBestFitHeap` 堆模拟，而不是闭式对齐公式）也在那里记录。把两者都当作继承事实，而不是在本页重新推导。
---
## 相对 post_layout 的配置差异
### 目的
`GetSchedulerConfig` (`0x10974aa0`) 是单个已编译函数。无论阶段如何，它都会构建同一个 137 字节 `SchedulerConfig` POD；融合前调用与 post_layout 调用的差异在于**调用点传入的参数**。POD 中有三个字段由融合前位置会以不同方式提供的输入驱动。
### 三个差异
```c
// GetSchedulerConfig(out, mem_limit_bytes=a2, env=a3, target=a4, is_lem=a5) // 0x10974aa0
char v45 = *(BYTE*)(env + 4097); // xla_tpu_rerun_lhs_post_sc_assignment (line 29)
// (A) memory_limit @ POD off 80 (lines 66-71)
out.memory_limit = -1;
if (v45) out.memory_limit = mem_limit_bytes;
if (EnableSchedulerMemoryPressureTracking(...)) out.memory_limit = mem_limit_bytes;
// (B) schedule_send_recvs @ off 120 (line 76)
out.schedule_send_recvs = v45 | ScheduleSendRecvs(...);
// (C) enable_send_recv_aggregation @ off 127 (line 90)
out.enable_send_recv_aggregation = env[4789] | v45;| POD 字段（偏移） | post_layout（可运行） | post_layout_pre_fusion（此变体） |
|---|---|---|
v45 来源 (env+4097) | 在 SC-assignment 重跑时可能为 1 | 始终为 0（不是重跑） |
memory_limit (80) | -1，或在压力跟踪或重跑时为 mem_limit_bytes | -1，除非压力跟踪（重跑路径不可达） |
schedule_send_recvs (120) | v45 | ScheduleSendRecvs(...) — 重跑时强制开启 | 仅 ScheduleSendRecvs(...) |
enable_send_recv_aggregation (127) | env[4789] | v45 — 重跑时强制开启 | 仅 env[4789] |
mem_limit_bytes 参数 | 最终融合后预算 | 融合前预算（融合缓冲区更少） |
| reserved-SparseCore 集合 | 从 GetReservedSparseCores 填充 | 空（SC assignment 融合前尚未运行） |
POD 中其他所有内容都相同：按 kind 的重叠限制（偏移 0-72，例如 collective_permute_overlap_limit = env[3964]，以及在 EnableSparseCoreOffloadQueuingInLhs 下为 INT32_MAX 的 all_reduce_overlap_limit）、偏移 123 的 use_real_cost_model 字节、偏移 122-125 的打包 bool 四元组，以及偏移 134-135 的 disable_loop_consume_resource / enable_idle_count_extension 字节。完整 137 字节 schema 在 post_layout 页面列成表格；此变体不改变任何偏移。
陷阱 —
v45字节和 reserved-SparseCore 集合是从结构上区分融合前配置与 post_layout 配置的唯一内容，并且两者在融合前都被强制为“首次运行 / 没有保留项”状态。重新实现者若把 post_layout 配置原样复制到融合前槽位，会错误地携带已填充的 reserved-SC 集合（SC assignment 之前不存在的信息）和可能置位的v45（融合前没有意义的重跑语义）。两者都要清空。
上下文差异 — 估算器与 tracker
目的
AddPass<LatencyHidingScheduler> 接受一个 SchedulingContext（估算器 + tracker）和一个 DefaultSchedulerCore。两者都由共享 helper 构建；融合前差异仍在输入中，主要是传给 tracker 的空 reserved-SparseCore 集合。
LatencyEstimator — 相同选择
GetLatencyEstimator (0x10974e00) 根据 SchedulerConfig 偏移 123 (use_real_cost_model) 分支：
c
function GetLatencyEstimator(target, config, env): // 0x10974e00
if config.use_real_cost_model != 1: // off 123, line 55
if PGLE_profile_path set:
FATAL "Attempted to use profile-guided latency estimator without the real "
"cost model. Either enable --xla_tpu_scheduler_using_real_cost_model "
"or unset --xla_tpu_impure_latency_hiding_scheduler_profile_path." // line 70
return ApproximateLatencyEstimator()
est = CostModelLatencyEstimator(target, config, env) // 0x10ff8a60, line 91
if PGLE_profile_path non-empty:
proto = tsl::ReadBinaryProto(path) // line 99
est = ProfileGuidedLatencyEstimator(est, proto) // line 104
return est
```text
这就是可运行 `post_layout` pass 使用的同一个 `GetLatencyEstimator`；融合前变体以完全相同的方式选择估算器（选择是 `use_real_cost_model`、配置字节和 PGLE 标志的函数，没有任何阶段依赖）。有关成本模型族，见 [`cost/overview`](../cost/overview.md)；有关 `CostModelLatencyEstimator` 路径，见 [`learned-cost-model-client`](../cost/learned-cost-model-client.md)。
### AsyncTracker — 空 reserved-SC 集合
`GetTpuAsyncTracker` (`0x10975520`) 包装 `TpuAsyncTracker::Create` (`0x10ffb3e0`，第 179 行调用)，并把 `absl::flat_hash_set<int64_t>` reserved-SparseCore 集合作为最后一个参数。在可运行 `post_layout` 路径中，该集合由 `GetReservedSparseCores` 填充（`RunHloScheduler` 第 873 行），也就是从重叠池中保留出来的 SparseCore ID，而这些信息*只在* SparseCore assignment 之后存在。融合前调用发生在 SC assignment 之前，因此该集合是**空的**。基础资源分类（ICI link、HBM bank、megacore queue、host-DMA tap，以及 `EnableSparseCoreOffloadQueuingInLhs` 下的 SparseCore offload queue）不变；[`scheduler-resourcetype-model`](scheduler-resourcetype-model.md) 页面记录了它。空 reserved-SC 集合是此变体唯一的 AsyncTracker 差异。
### DefaultSchedulerCore — 相同接线
core ctor (`0x10976ce0`) 通过 `vmovups` 块将 137 字节配置 POD 从 `[rcx]…[rcx+0x69]` 复制到 `[rdi+0x10]…[rdi+0x79]`（第 23-32 行），并安装四个 `std::function` 槽。在可运行构建中，`candidate_compare_` 和 `post_step_mutator_` 是**空的**（`RunHloScheduler` 第 1000 / 1002 行），因此 core 回退到内置的 `ReadySetLt` 启发式；`should_schedule_` 是由 `EnableIlpLatencyHidingScheduler` 选择的 `$_1`（ILP）或 `$_2`（常规）分类器（第 1084-1107 行），`target_overlap_limit` 是 `RunTensorCoreAsyncOpScheduler::$_0`（第 1040 行）。融合前构建会以相同方式接线；分类器切换是 [ILP variant](lhs-ilp-variant.md)，与阶段正交。
---
## post_layout 与 post_layout_pre_fusion 概览
| 方面 | post_layout（可运行） | post_layout_pre_fusion（本页） |
|---|---|---|
| 预期阶段 | 第 7 阶段，`RunHloScheduler` 中的 `async_scheduling` | 第 6 阶段，`"Pre main fusion"` 流水线 |
| 在 v0.0.40 中已接线？ | **是**（`RunHloScheduler` 第 1137 / 1411 行） | **否** — 融合前流水线中没有 `AddPass<LHS>` |
| 能看到融合结果？ | 是（在 `"Main fusion"` 之后运行） | 否（本应在融合之前运行，这是其意图） |
| 有基础 schedule？ | 有（`HloMemorySchedulerWithBrkgaFallback`，第 567 行） | **无** → 若添加则 `has_schedule()` FATAL |
| `RunImpl` | `0x136321a0`（共享） | `0x136321a0`（共享，主体相同） |
| `SchedulerConfig` 构建器 | `GetSchedulerConfig` (`0x10974aa0`) | `GetSchedulerConfig`（相同） |
| `v45` 重跑字节 (`env+4097`) | 可能为 1（SC-assignment 重跑） | 始终为 0 |
| `memory_limit` (off 80) | `-1`，除非压力跟踪**或** `v45` | `-1`，除非压力跟踪 |
| `mem_limit_bytes` 参数 | 最终融合后预算 | 融合前预算 |
| reserved-SparseCore 集合 | 来自 `GetReservedSparseCores` | 空（SC assignment 尚未运行） |
| LatencyEstimator | `GetLatencyEstimator`（相同） | `GetLatencyEstimator`（相同） |
| `candidate_compare_` / mutator | 空（启发式） | 空（启发式） |
| 异步分类器 | `$_2`，或 ILP gate 下的 `$_1` | `$_2`，或 ILP gate 下的 `$_1` |
v0.0.40 中唯一可运行的 LHS 是 `post_layout`。此变体是为死掉的融合前槽位配置的同一个 `RunImpl`；记录它就是记录 (a) 槽位为何为空，以及 (b) 未来构建需要的精确三项输入差异。
---
## 函数映射
| 函数 | 地址 | 作用 |
|---|---|---|
| `HloOptimizeAfterLayoutAssignment` | `0x10962660` | 第 6 阶段驱动器；构建 `"Pre main fusion"` 流水线（第 1169 行）— 未添加 LHS |
| `PostOptimizationPipeline` | `0x1093fd40` | 第 7 阶段驱动器；调用 `RunHloScheduler`（从 `RunHloPasses` 第 1987 行到达） |
| `RunHloScheduler` | `0x1096fac0` | 唯一的 LHS `AddPass` 位置（1137 / 1411）；由 `PostOptimizationPipeline` 构建 |
| `LatencyHidingScheduler::RunImpl` | `0x136321a0` | 共享 pass 主体；第 280 行 `has_schedule()` FATAL，重试循环 |
| `GetSchedulerConfig` | `0x10974aa0` | 137 字节 `SchedulerConfig`；第 29 行 `v45 = env[4097]` |
| `GetLatencyEstimator` | `0x10974e00` | Approximate / CostModel / PGLE 选择（字节 123） |
| `GetTpuAsyncTracker` | `0x10975520` | 包装 `TpuAsyncTracker::Create`；接受 reserved-SC `flat_hash_set` |
| `TpuAsyncTracker::Create` | `0x10ffb3e0` | TPU 资源模型 + reserved-SC 集合（第 179 行调用点） |
| `DefaultSchedulerCore` ctor | `0x10976ce0` | 复制配置 POD (`vmovups`)；安装四个 `std::function` hook |
| `AddPass<LatencyHidingScheduler, …>` | `0x10975c40` | 模板化 pass 添加器 — 只从 `RunHloScheduler` 引用 |
| `HloMemorySchedulerWithBrkgaFallback` ctor | (`0x1096fac0` 第 567 行) | 第 7 阶段基础内存 schedule（`has_schedule()` 前置条件） |
| `IsPassDisabled` | `0x12fd8340` | `"latency-hiding-scheduler"` / `"async-op-scheduler"` gate |
| `EnableSchedulerMemoryPressureTracking` | `0x1d6b66e0` | `memory_limit` 三态 |
| `EnableIlpLatencyHidingScheduler` | `0x1d6b7e00` | `$_1`/`$_2` 分类器 gate（ILP 变体） |
---
## 置信度摘要
| 结论 | 证据 |
|---|---|
| 融合前槽位没有 `AddPass<LatencyHidingScheduler>` | `0x10962660` — 只有 `"Pre main fusion"`（1169）/ `"Main fusion"`（2279）；没有 LHS / 没有 `RunHloScheduler` 引用 |
| 规范 `AddPass<LHS>` 模板只从 `RunHloScheduler` 引用 | 对 `AddPassINS_22LatencyHidingScheduler` 执行 `rg --no-ignore` → 1 个文件（其自身主体） |
| 两个可运行 LHS 位置都是第 7 阶段（融合后） | `RunHloScheduler` 第 1137 / 1411 行 |
| 槽位失效，因为 `RunImpl` 在 `!has_schedule()` 上 FATAL | `0x136321a0` 第 280 行 `LogMessageFatal` + 前置条件字符串 |
| 基础 schedule 只存在于第 7 阶段 | `RunHloScheduler` 第 567 行的 `HloMemorySchedulerWithBrkgaFallback` |
| `v45 = env[4097]`；驱动 `memory_limit`/`schedule_send_recvs`/aggregation | `GetSchedulerConfig` 第 29 / 66-71 / 76 / 90 行 |
| 融合前配置差异 = `v45==0`、融合前预算、空 reserved-SC | 根据 `GetSchedulerConfig` + `GetReservedSparseCores`（第 873 行）阶段顺序推导 |
| `next_memory_limit` 松弛因子 `qword_A2DFD10` | `RunImpl` 第 892 / 918 行 `vmulsd`（十进制 `0.9` 在 core 页面解析） |
| 空 reserved-SC 集合是 AsyncTracker 差异 | `GetTpuAsyncTracker` 最后的 `flat_hash_set<long>` 参数；SC assignment 晚于融合前 |
| `GetLatencyEstimator` 分支 + PGLE FATAL 相同 | `0x10974e00` 第 55 / 70 / 91 / 104 行 |
---
## 交叉引用
- [LatencyHidingScheduler 核心](latency-hiding-scheduler-core.md) — `RunImpl` 驱动的共享 `ScheduleComputation` drain loop、22 键比较器和 `TpuAsyncTracker` 资源模型。
- [LHS post_layout](lhs-post-layout.md) — 可运行的第 7 阶段 LHS；本页是它的差异说明。完整 137 字节 `SchedulerConfig` schema、双流水线 `RunHloScheduler` 结构和 reserved-SparseCore 集合都记录在那里。
- [LHS ILP 变体](lhs-ilp-variant.md) — `EnableIlpLatencyHidingScheduler` `$_1`/`$_2` 分类器切换，与调度阶段正交。
- [Scheduler ResourceType 模型](scheduler-resourcetype-model.md) — `AsyncTracker` 强制执行的 ICI/HBM/megacore/host-DMA/SparseCore 资源分类。
- [调度概览](overview.md) — 调度在 lowering 与 encoding 之间的位置。
- [成本模型概览](../cost/overview.md) — 供给 `GetLatencyEstimator` 的延迟/成本模型族。
- [Learned Cost Model Client](../cost/learned-cost-model-client.md) — `CostModelLatencyEstimator` 真实成本模型路径。
- **二进制：** `extracted/libtpu-0.0.40-cp314-cp314-manylinux_2_31_x86_64/libtpu/libtpu.so`（build-id `89edbbe81c5b328a958fe628a9f2207d`）
- **索引条目：** 第 VIII 部分 — 指令调度与 Bundle Packing — [返回索引](../index.md)