编译阶段
地址、构建 ID 和符号名称适用于
libtpu-0.0.40-cp314轮中的libtpu.so(构建 ID89edbbe81c5b328a958fe628a9f2207d)。其他版本有所不同;将每个 VA 视为版本固定。
摘要
TPU 程序是通过端到端运行 有序相序 来编译的:前端 HLO 预先清理并规范化传入的图形,然后分片/SPMD 分区按设备专门化它,然后布局分配修复物理张量布局,然后布局后 HLO 细化和融合塑造图形,然后 HLO 被交给 MLIR 进行方言下降(StableHLO → MHLO → TLP → tpu 方言 → LLO),然后是内存空间分配和调度价格并订购 LLO,然后捆绑打包器将 ops 装入 VLIW 字中,最后将程序序列化为 TpuExecutable。该页面拥有顶级有序阶段列表和每个阶段的入口符号 - 编译器每个模块遍历一次的脊柱。
该脊柱有两种观点,并不矛盾。 多阶段单独编译视图公开了五个公共 C++ 阶段函数 - xla::CompilePhase0StablehloToHlo 到 xla::CompilePhase3Linking - 每个函数消耗并生成序列化的 PjRtPartialProgramProto 跨度,以便调用者可以在阶段之间进行检查点。 整体视图在一个 DeepseaCompilerBase::RunHloPasses 调用中运行相同的工作,该调用直接构建并运行 HLO xla::HloPassPipeline 容器。单独编译 API (0/1/2a/2b/3) 中的阶段编号比内部管道构建器阶段(PreOptimization → sharding → SPMD-prep → through-layout → LayoutAssignment → PostOptimization → MLIR)更粗。该页面对两者进行了协调并给出了每个的条目符号。
这是订购页面。固定了各阶段的顺序和切入点;它不会重新派生 IR 层堆栈(HLO/MHLO/StableHLO/CHLO/VHLO/TLP/tpu/LLO 操作计数和转换传递)——该概念存在于 概述.md 上并且是链接的,而不是重复的。完整枚举的 HLO 预传递表位于 hlo-pre-passes.md 上,每类传递目录位于 hlo-pass-registry.md 上;此页面仅记录设置为排序头部的 front-of-pipeline 预通道。
对于重新实现,相序契约为:
- **阶段顺序是总的,每个阶段都会消耗前一阶段的输出。**阶段 1–5(在管道构建器编号中)是 HLO→HLO;到 MLIR 的移交是第 6 阶段。没有任何阶段会乱序运行,也没有任何阶段会被整体跳过 - 阶段内的各个“通道”都是标记门控的,但阶段脚手架始终运行。
xla::HloPassPipeline是唯一的 HLO 容器类型。 没有 TPU 专用管道类;每个 TPU 通行证均源自xla::HloPassInterface,并添加AddPass<T>。不变检查器通过单独的AddInvariantChecker<T>API 添加,并在管道中的“每次”通过后重新验证,而不是一次。- HLO 预传递是排序的前面。
PreOptimizationPipeline首先运行(输入清理:调度注释合法化、自定义调用扩展器、动态形状填充、精度重写器),然后再进行任何分片、布局或融合决策。 - MLIR/LLO/MSA/调度/捆绑打包阶段在所有 HLO 工作的下游运行。 布局分配、MSA 和两个调度程序都需要完全预传递、分片、融合的 HLO 模块作为先决条件;捆绑打包程序需要降低 LLO。
- 序列化关闭序列。 编译后的程序通过
TpuExecutable_SerializeC-ABI 导出作为TpuExecutable发出;单独编译流程将每个相边界重新序列化为PjRtPartialProgramProto。
| 顶级编排器(整体) | xla::jellyfish::DeepseaCompilerBase::RunHloPasses(Target const&, RunHloPassesConfig const&, …) @ 0x1093a420 — 管道建设过载; xla::Compiler 虚拟入口 RunHloPasses(unique_ptr<HloModule>, StreamExecutor*, CompileOptions const&) @ 0x109152a0 代表在这里(参见 概述.md) |
| 单独编译阶段fns | Phase0 @ 0xf84de60 · Phase1 @ 0xf84ee00 · Phase2a @ 0xf850840 · Phase2b @ 0xf852180 · Phase3 @ 0xf852f40 |
| 管道建设者 | CreateHloPipeline @ 0x1093efe0 · PreOptimizationPipeline(在 0x10948840 AddPass thunk 中) · HloOptimizeThroughLayoutAssignment @ 0x1094ad80 · PostOptimizationPipeline @ 0x1093fd40 · MaybeAddInvariantCheckers @ 0x10944600 |
| 分片/SPMD | AddAutoShardingAndRelatedPasses @ 0x10939c40 · AddTpuPartitioningPasses @ 0x1278a440 · TpuSpmdPartitioner (AddPass 0x1278b8a0) |
| 验收测试 | xla::TpuHloSupportChecker::RunImpl @ 0x11071480 (name() @ 0x11071780,返回 "tpu_hlo_support_checker") |
| 布局 | Phase2PreLayoutAssignment @ 0x1094e0e0 · TpuLayoutAssignment(xla::LayoutAssignment 的子类) |
| MLIR 切换 | mlir::mhlo::HloLegalizeToStablehloPass::runOnOperation @ 0x16adcea0;最终预 MLIR HLO 通过 ConvertFrontendAttributesToBackendConfig::RunImpl @ 0x10a148a0 |
| MSA | xla::jellyfish::RunMemorySpaceAssignment @ 0x12fc3080 |
| 调度 | (anon)::RunHloScheduler @ 0x1096fac0 |
| 捆绑包装 | GlobalBundlePacker::PackInstruction @ 0x10a875a0(参见 LLO捆绑包装) |
| 序列化 | TpuExecutable_Serialize @ 0xeabea80 · TpuExecutableSerialize_WriteToArray @ 0xeabeba0 |
| 信心 | 已确认(字节锚定),除非行或标注另有说明 |
两种编号方案
libtpu 通过两个不同的表面公开编译管道,重新实现必须使它们保持一致,因为“阶段”一词在每个表面中意味着不同的东西。
单独编译阶段(公共 C++ API)。 从 RTTI 恢复,五个自由函数各自采用 (CompileOptions, absl::Span<const PjRtPartialProgramProto>, const PjRtTopologyDescription&) 并返回部分程序的 StatusOr。它们是 PJRT 单独编译流程的可检查点边界 - 调用者可以运行阶段 0,保留中间原型,并稍后在阶段 1 恢复。 Phase1HloOptimizations 主体证实了这种形状:它以 HloModuleProto / PjRtPartialProgramProto 粒度工作,反序列化 HLO 模块原型,运行 HLO 优化阶段,然后重新序列化。
| Sep-comp阶段 | 符号 | VA | 它涵盖什么 |
|---|---|---|---|
| 0 | xla::CompilePhase0StablehloToHlo | 0xf84de60 | StableHLO/MHLO 前端导入 → HLO |
| 1 | xla::CompilePhase1HloOptimizations | 0xf84ee00 | 整个 HLO 传递管道(预传递 → 分片 → 布局 → 后选择) — 与 RunHloPasses 相同的工作 |
| 2a | xla::CompilePhase2aTlpLowering | 0xf850840 | HLO → MHLO → TLP MLIR 降低 |
| 2b | xla::CompilePhase2bDedupedLowering | 0xf852180 | 去重降低(tpu 方言发射之前/内的共享子图去重) |
| 3 | xla::CompilePhase3Linking | 0xf852f40 | LLO 链接,捆绑包,序列化为 TpuExecutable |
管道构建器阶段(内部)。 在 Phase1/RunHloPasses 内部,HLO 工作由管道构建器函数分解为更精细编号的阶段。这些是 hlo-pre-passes.md 订购表使用的阶段(其阶段 1-6)。映射为: sep-comp Phase 0 ≈ builder stage 0(前端导入); sep-comp 第 1 阶段 = 构建器阶段 1–5(所有 HLO→HLO); sep-comp 阶段 2a/2b = 构建者阶段 6(MLIR 血统)加上 LLO 级 MSA/计划通行证; sep-comp 第 3 阶段 = 捆绑包 + 链接 + 序列化。
**明白了——“第 2b 阶段重复数据删除降低”不是逆向工程,只是命名。**其重复数据删除算法(其折叠的相同子图,以及在哪个 IR 级别)未恢复。仅将上面的行视为相边界。 [置信度:确认该相存在@
0xf852180;其内部机制低。]
有序相表
编译器对每个模块执行以下序列一次。下图是从 RunHloPasses (0x1093a420) 及其被调用者恢复的规范排序;每个方框都是一个阶段,其输入符号位于右侧。阶段1-5是纯HLO→HLO;第 6 阶段将 HLO 交给 MLIR;阶段 7-10 在降低后运行。
StableHLO / MHLO module from PJRT frontend
│ CompilePhase0StablehloToHlo 0xf84de60
▼
HLO module
│
▼ (all of the below = CompilePhase1HloOptimizations 0xf84ee00 / RunHloPasses 0x1093a420)
┌───────────────────────────────────────────────────────────────────────────┐
│ PHASE 1 — HLO PRE-PASSES (input scrub) PreOptimizationPipeline │
│ front pass: HloDomainIsolator (sharding-domain bracketing) │
│ custom-call expanders: Tpu{Cholesky,Qr,Eigh,TriangularSolve}Expander, │
│ FftExpander, LuDecompositionExpander │
│ dynamic shape: DynamicIndexSplitter, DynamicDimensionSimplifier, │
│ DynamicPadder │
│ precision/dtype: TpuInt2AutoUpDownCaster, XPrecisionRewriter (x6/x9/x128), │
│ TpuHloPrecisionTracer, BitcastDtypesExpander │
│ misc: BatchNormExpander, TpuCallInliner(MustFuse), HloDCE │
└───────────────────────────────────────────────────────────────────────────┘
│ (CreateHloPipeline 0x1093efe0 opens the main pipeline)
▼
┌───────────────────────────────────────────────────────────────────────────┐
│ PHASE 2 — SHARDING AddAutoShardingAndRelatedPasses│ 0x10939c40
│ manual flow : ShardingPropagation │
│ auto flow : TpuAutoSharding → ShardingPropagation │
│ Shardy flow : sdy::ShardyXLA (when JAX/Shardy frontend detected) │
│ prep : FlattenCallGraph, TpuCallInliner, TupleSimplifier │
└───────────────────────────────────────────────────────────────────────────┘
│
▼
┌───────────────────────────────────────────────────────────────────────────┐
│ PHASE 3 — SPMD PARTITIONING PREP AddTpuPartitioningPasses │ 0x1278a440
│ SpmdPrepare, HloConstantSplitter, TpuPartitionAssignment, │
│ TpuSpmdConcatRewriter, TpuSpmdPartitioner (per-partition HLO + collectives)│
└───────────────────────────────────────────────────────────────────────────┘
│
▼
┌───────────────────────────────────────────────────────────────────────────┐
│ PHASE 4 — THROUGH-LAYOUT HLO OPT HloOptimizeThroughLayoutAssignment │ 0x1094ad80
│ acceptance : TpuHloSupportChecker (the canonical TPU HLO filter) │ 0x11071480
│ simplify : TpuAlgebraicSimplifier, HloCSE, HloConstantFolding │
│ collective : AllReduceSimplifier, TpuAllGatherSimplifier, AllToAllDecomposer│
│ shape/conv : ConvolutionFolding, TransposeFolding, Tpu*Rewriter family │
│ gather/scatter: TpuGatherScatterFlattener / TpuGather/ScatterExpander │
│ ───────────────────────────────────────────────────────────────────── │
│ LAYOUT ASSIGNMENT Phase2PreLayoutAssignment / TpuLayoutAssignment │ 0x1094e0e0
└───────────────────────────────────────────────────────────────────────────┘
│
▼
┌───────────────────────────────────────────────────────────────────────────┐
│ PHASE 5 — POST-LAYOUT HLO REFINEMENT PostOptimizationPipeline │ 0x1093fd40
│ HloCSE (post-layout), AddOriginalValue (provenance metadata), │
│ FUSION (main fusion runs here — see fusion-patterns.md), │
│ final pre-MLIR pass: ConvertFrontendAttributesToBackendConfig │ 0x10a148a0
└───────────────────────────────────────────────────────────────────────────┘
│ CompilePhase2aTlpLowering 0xf850840 / CompilePhase2bDedupedLowering 0xf852180
▼
┌───────────────────────────────────────────────────────────────────────────┐
│ PHASE 6 — MLIR DIALECT DESCENT │
│ HloLegalizeToStablehloPass (HLO → StableHLO bytecode) │ 0x16adcea0
│ StableHLO → MHLO → TLP → tpu MLIR dialect → LLO │
│ (IR-layer stack & op counts: see overview.md) │
└───────────────────────────────────────────────────────────────────────────┘
│
▼
┌───────────────────────────────────────────────────────────────────────────┐
│ PHASE 7 — MEMORY-SPACE ASSIGNMENT RunMemorySpaceAssignment │ 0x12fc3080
│ (HBM/VMEM/SMEM placement; LSRAv2 live-range alloc — see msa-overview.md) │
└───────────────────────────────────────────────────────────────────────────┘
│
▼
┌───────────────────────────────────────────────────────────────────────────┐
│ PHASE 8 — SCHEDULING RunHloScheduler │ 0x1096fac0
│ base order (HloMemorySchedulerWithBrkgaFallback) → async_scheduling │
│ pipeline (LatencyHidingScheduler) + LLO MXU/MRB (see sched/overview.md) │
└───────────────────────────────────────────────────────────────────────────┘
│
▼
┌───────────────────────────────────────────────────────────────────────────┐
│ PHASE 9 — LLO BUNDLE PACKING GlobalBundlePacker::PackInstr. │ 0x10a875a0
│ forward greedy earliest-legal-bundle VLIW packing │
└───────────────────────────────────────────────────────────────────────────┘
│ CompilePhase3Linking 0xf852f40
▼
┌───────────────────────────────────────────────────────────────────────────┐
│ PHASE 10 — SERIALIZATION / LINK TpuExecutable_Serialize │ 0xeabea80
│ link LLO program → TpuExecutable → WriteToArray │ 0xeabeba0
└───────────────────────────────────────────────────────────────────────────┘
│
▼
TpuExecutable bytes
```text
排序是由数据依赖性强制的,而不是约定。分片(第 2 阶段)必须先于 SPMD 分区(第 3 阶段),因为分区程序会在每个操作上使用分片注释。布局分配(第 4 阶段内)必须先于融合(第 5 阶段),因为融合合法性取决于物理图块布局。 MLIR 下降(第 6 阶段)必须遵循所有 HLO 工作,因为 `tpu` 方言没有未预先通过的 HLO 图的概念。 MSA(阶段 7)必须先于调度(阶段 8),因为调度程序针对分配的内存空间按周期对溢出/重新填充进行定价。捆绑打包(第 9 阶段)是优化阶段中的最后一个,因为槽合法性是每个 VLIW 字的微观决策,仅在程序为 LLO 时才存在。
---
## 第 1 阶段 — HLO Pre-Pass 前端
订购的开头是 **`PreOptimizationPipeline`** — 内置在 `deepsea_compiler_hlo_passes.cc` 的 `(anonymous namespace)` 内部的 `xla::HloPassPipeline`。这里的每一次传递都会消耗 HLO 并发出 HLO;还没有任何东西被分割、布置或降低。这是“唯一”将自由形式前端图规范化为编译器其余部分可以采用的形状的地方。
恢复的管道前端通道集(`name()` 锚定的、管道确认的子集 — 具有 HLO 输入/输出不变量的完整枚举位于 [hlo-pre-passes.md](hlo-pre-passes.md) 上):
| 订单 | 合格等级 | 角色 |
|---|---|---|
| 前部 | `xla::HloDomainIsolator`(谓词 `$_7`) | 在显式 `kDomain` 操作中包装分片区域(仅当 `EnableDomainPasses()` 时) |
| 早 | `xla::DynamicIndexSplitter` | 将多维动态索引拆分为标量 |
| 早 | `xla::BatchNormExpander(true,true,true)` | 降低原始算术的批量范数 |
| 扩展器 | `xla::TpuCholeskyExpander` / `TpuQrExpander` / `TpuEighExpander` / `TpuTriangularSolveExpander` | 将 linalg 自定义调用分解为点/三角求解图 |
| 扩展器 | `xla::FftExpander(Target const&)` / `xla::LuDecompositionExpander` | FFT(TPU-aware radix-2)/LU分解 |
| rng | `xla::jellyfish::TpuRngBitGeneratorExpander` (+ `TpuRngBitGeneratorTupleDecomposer`) | Philox/ThreeFry;元组 `(state,output)` |
| dtype | `xla::jellyfish::TpuInt2AutoUpDownCaster` | 括号 int2 算术与 `Convert`↔int8(MXU 线最小值为 int8) |
| 内联 | `xla::jellyfish::TpuCallInliner(MustFuseInlineMode)` | 内联 `must_fuse` 标记的早期被调用者 |
| 融合准备 | `xla::jellyfish::UserGuidedFusionIdAssigner` | 将 `frontend_attribute: fusion_id` 字符串转换为整数后端配置 |
| 动态 | `xla::DynamicDimensionSimplifier` → `xla::DynamicPadder(...)` | 折叠然后静态垫动态形状 |
| 精密 | `xla::jellyfish::XPrecisionRewriter(kX128Precision)` 然后 `XPrecisionRewriter()` | x128→8步,x6/x9→2/3步点累积链 |
| 精密 | `xla::jellyfish::TpuHloPrecisionTracer` / `xla::BitcastDtypesExpander` | 填写缺失的`precision_config`;扩展仅 dtype 的位广播 |
| 清理 | `xla::HloDCE`(在大多数阶段之间重新运行) | 死指令/计算 |
| 边界 | `xla::HloDomainRemover("sharding", ...)` | 剥离 `kDomain`,保留分片作为属性 — 始终在阶段 2 之前运行 |
> **注意 — 定点循环。** 多个通道包含在 `xla::HloPassFix<P>` 中并重新运行以收敛。 `RunToFixPoint` 存在于 `HloPassFix<xla::ReduceWindowRewriter>` (@ `0x14bd0980`)、`HloPassFix<xla::jellyfish::TpuReduceWindowRewriter>` (@ `0x109589e0`) 和 `HloPassFix<xla::HloDCE>` (@ `0x1d6d7a60`) 的二进制文件中。迭代限制和不收敛崩溃由存在字符串的真实标志控制:`xla_tpu_hlo_pass_fix_pipeline_iteration_limit`(每个管道上限)、`xla_unsupported_crash_on_hlo_pass_fix_max_iterations`(如果达到上限则中止)和 `xla_hlo_pass_fix_detect_cycles`。 `HloPassFix`定点机制和这些标志字符串已确认;未在标志接线之外跟踪大写命中的确切运行时行为。
>
> **发现 - `TpuHloSupportChecker` 是接受门,它在第 4 阶段运行,而不是第 1 阶段。** 想要尽早拒绝不支持的操作的重新实现会试图将支持检查放在前面。 libtpu 不:`TpuHloSupportChecker::RunImpl` (`0x11071480`) 在 `HloOptimizeThroughLayoutAssignment`*之后*分片和 SPMD 准备内部运行,因为预传递合法地*引入*操作(扩展器输出、分区集合),这些操作本身必须通过检查。在扩展之前进行检查会拒绝实际可编译的程序。检查器永远不会改变模块 - 它会遍历每个 `HloComputation` 并使用 `ShapeUtil::ValidateShapeWithOptionalLayout` 验证每个结果 `Shape`,在第一个不受支持的形状上返回错误 `Status`。
---
## 阶段 2-3 — 分片和 SPMD 分区
两个构建器函数拥有该区域。 `AddAutoShardingAndRelatedPasses` (`0x10939c40`) 决定程序如何分片; `AddTpuPartitioningPasses` (`0x1278a440`) 通过将单设备图重写为带有集合的每分区图来*应用*分片。
三种分片流共存,通过flags和前端属性检测进行选择:
- **手动** — 预先存在的分片注释由 `xla::ShardingPropagation` 传播,直到每个操作都有分片(或 `Replicated`)。
- **自动** - `xla::TpuAutoSharding`(封装开源 `xla::AutoSharding`)合成注释,然后 `ShardingPropagation` 填充其余部分。由 `xla_tpu_auto_sharding_enabled` 标志系列控制。
- **Shardy** — `xla::sdy::ShardyXLA` 在检测到 JAX/Shardy 前端时运行(其 `AddPass` thunk 出现在 `0x10959e80`、`0x1278b3a0`、`0x14bbc200` 中),将 Shardy 格式的分片操作转换为 HLO分片注释。
分区器 — `xla::jellyfish::TpuSpmdPartitioner`(AddPass thunk `0x1278b8a0`)是开源 `SpmdPartitioner` 的 TPU 特定子类 — 然后在 `TpuPartitionAssignment`(设备 ID 元数据)的协助下生成具有显式集合的每个分区 HLO `TpuSpmdConcatRewriter`。
> **注意 - 选择器逻辑未恢复。** 三个流程中的哪一个运行由标志加上构建器内部更深入的前端属性检测决定;确切的选择器分支没有被隔离。 [信心:确认所有三个流程都存在并且由这些构建者添加;精确调度条件为低。] 请参阅 [自动分片-spmd.md](auto-sharding-spmd.md) 和 [分片传播.md](sharding-propagation.md)。
---
## 阶段 4-5 — 布局分配和布局后细化
`HloOptimizeThroughLayoutAssignment` (`0x1094ad80`) 运行大部分 HLO 优化 - 代数简化(`TpuAlgebraicSimplifier`,开源 `AlgebraicSimplifier` 的超集)、CSE、常量折叠、集体简化、卷积折叠和 TPU 形状重写器系列(`TpuBroadcastRewriter`、`TpuDegenerateDimensionRewriter`、`TpuReduceRewriter`、`HloPassFix<TpuReduceWindowRewriter>`)—然后交给**布局分配**(`Phase2PreLayoutAssignment` @ `0x1094e0e0` 驱动 `TpuLayoutAssignment`,`xla::LayoutAssignment` 的子类)。请参阅 [布局分配.md](layout-assignment.md)。
`PostOptimizationPipeline` (`0x1093fd40`) 然后细化现在布局的图:布局后 `HloCSE(false)`、通过 `xla::AddOriginalValue` 进行来源标记、**主要融合**(记录在 [fusion-patterns.md](fusion-patterns.md) 和 [fusion-cost-model.md](fusion-cost-model.md) 上),最后是 `xla::jellyfish::ConvertFrontendAttributesToBackendConfig::RunImpl` (`0x10a148a0`) — **最后一个 HLO 域通道**。它将已知的 TPU `frontend_attribute` 密钥解析为类型化的 `backend_config` protobuf,而未知的密钥保持不变,因此下游 MLIR 导入看不到任何 TPU 私有的内容。它处理的特定键并未全部恢复; `fusion_id` 属性(之前由 `UserGuidedFusionIdAssigner` 分配,其源路径存在 `user_guided_fusion_id_assigner.cc`)和 `scheduling_group_id` 是二进制文件中确认的两个属性字符串。 [置信度:确认该pass存在@`0x10a148a0`并转换前端属性;整个按键设置较低。]
> **注意 - 不变的检查器连续重新运行。** `MaybeAddInvariantCheckers` (`0x10944600`) 在每个的头部重新添加 `HloVerifier` (+ `TpuVerifierMetadata`)、`LegalizeSchedulingAnnotations`(作为检查器)和 `HloCycleDetection`通过 `AddInvariantChecker<T>` 嵌套管道,因此它们在每次通过后重新验证而不是一次。仅在阶段边界运行验证的重新实现将接受 libtpu 在中间阶段拒绝的格式错误的中间体。 [信心:已确认。]
---
## 第 6 阶段 — MLIR 方言下降
`ConvertFrontendAttributesToBackendConfig` 运行后,HLO 模块将被交给 MLIR。第一个 MLIR 通道是 `mlir::mhlo::HloLegalizeToStablehloPass` (`runOnOperation` @ `0x16adcea0`),它将 HLO 转换为 StableHLO 字节码;从那里继续下降 StableHLO → MHLO → TLP → `tpu` 方言 → LLO。本页**不**枚举方言操作数、每种方言的转换通道或合法化器内部结构 - 即 [概述.md](overview.md) 上的 IR 层堆栈材料,以及 [tpu-to-llo-ods.md](tpu-to-llo-ods.md)、[方言-转换-legalizer.md](dialect-conversion-legalizer.md) 上的降低合法化器以及 [马赛克-overview.md](mosaic-overview.md) 和 [马赛克-overview.md](mosaic-overview.md) 上的马赛克/自定义调用路径[自定义呼叫降低.md](custom-call-lowering.md)。在单独编译编号中,此血统是 `CompilePhase2aTlpLowering` (`0xf850840`),后面是 `CompilePhase2bDedupedLowering` (`0xf852180`)。
---
## 阶段 7–10 — MSA、调度、捆绑打包、序列化
这些阶段在降级后运行,属于其他页面所有;此页面仅修复它们在排序中的位置及其输入符号。
- **阶段 7 — 内存空间分配。** `xla::jellyfish::RunMemorySpaceAssignment` (`0x12fc3080`) 将张量跨 HBM / VMEM / SMEM 放置并运行实时范围分配 (LSRAv2)。它先于调度,因为调度程序根据分配的空间对溢出/重新填充进行定价。请参见 [msa-overview.md](msa-overview.md)、[msa-分配-segment.md](msa-allocate-segment.md)、[msa-reservation-hbm-policy.md](msa-reservation-hbm-policy.md)。
- **阶段 8 — 调度。** `(anon)::RunHloScheduler` (`0x1096fac0`) 运行双管道 HLO 调度程序 — 一个内存最小化基本顺序(`HloMemorySchedulerWithBrkgaFallback`,由 `hlo-memory-scheduler-with-brkga-fallback` 字符串锚定),然后运行 `async_scheduling` 管道( `LatencyHidingScheduler` 重叠重写)— 随后是 LLO 级 MXU 序列/MRB 分配。整个调度堆栈,包括第二个 LLO 级调度程序和为其定价的成本模型,都记录在 [../sched/overview.md](../sched/overview.md) 上。
- **阶段 9 — LLO 捆绑打包。** `GlobalBundlePacker::PackInstruction` (`0x10a875a0`) 是前向贪婪最早合法捆绑 VLIW 打包程序;内部硬件循环由单独的模调度器进行软件流水线处理。请参见 [../sched/llo-bundle-packing.md](../sched/llo-bundle-packing.md) 和 [../sched/bundle-modulo-scheduling.md](../sched/bundle-modulo-scheduling.md)。
- **阶段 10 — 序列化/链接。** `CompilePhase3Linking` (`0xf852f40`) 链接 LLO 程序并发出 `TpuExecutable`; C-ABI 导出 `TpuExecutable_Serialize` (`0xeabea80`) 和 `TpuExecutableSerialize_WriteToArray` (`0xeabeba0`) 写入字节。单独编译流程另外将每个相边界重新序列化为 `PjRtPartialProgramProto`。请参阅 [tpu-program-serialization.md](tpu-program-serialization.md)。
---
## 此页上没有的内容
- **IR 层堆栈**(HLO/StableHLO/MHLO/CHLO/VHLO/TLP/`tpu`/LLO 操作计数和方言到方言转换通道)— 请参阅 [概述.md](overview.md)。
- **完整的 HLO 预传递枚举**(所有 97 个管道提到的传递,以及每传递 HLO 不变量和 372 项 RTTI 传递目录) - 请参阅 [hlo-pre-passes.md](hlo-pre-passes.md) 和 [hlo-pass-registry.md](hlo-pass-registry.md)。
- **每遍转换算法**(每个重写器对图实际执行的操作) - 每个主题页面:[代数简化器.md](algebraic-simplifier.md)、[布局分配.md](layout-assignment.md)、[fusion-patterns.md](fusion-patterns.md) 等。
- **每 `TpuVersion` 管道分歧。** 反编译反映了按 `Target` / `TpuCompilationEnvironment` 标志分支的*一个*管道;哪些通道会被跳过,例如,TPU v3 与 v6e 没有隔离。 [信心:每代差异较低。]
---
## 交叉引用
- [概述.md](overview.md) — 编译器前端概述和 HLO → … → LLO IR 层堆栈(概念所有者)。
- [hlo-pre-passes.md](hlo-pre-passes.md) — 具有每遍 HLO 输入/输出不变量的完整有序 HLO 预传递表。
- [hlo-pass-registry.md](hlo-pass-registry.md) — 372 条目 `HloPassInterface` 类目录和 `name()` 字符串。
- [布局分配.md](layout-assignment.md) — `TpuLayoutAssignment`(第 4 阶段尾部)。
- [fusion-patterns.md](fusion-patterns.md) / [fusion-cost-model.md](fusion-cost-model.md) — 主要融合(第 5 阶段)。
- [自动分片-spmd.md](auto-sharding-spmd.md) / [分片传播.md](sharding-propagation.md) — 第 2-3 阶段详细信息。
- [msa-overview.md](msa-overview.md) — 第 7 阶段内存空间分配。
- [../sched/overview.md](../sched/overview.md) — 第 8 阶段调度堆栈(HLO LHS + LLO 捆绑包 + MXU/MRB)。
- [tpu-to-llo-ods.md](tpu-to-llo-ods.md) — `tpu` 方言 → LLO 降低(第 6 阶段尾部)。
- [tpu-program-serialization.md](tpu-program-serialization.md) — 第 10 阶段 `TpuExecutable` 序列化。
- [返回索引](../index.md)