内部 Pass 名称目录
本页中的所有地址和偏移都适用于
libtpu-0.0.40-cp314wheel 中的libtpu.so(libtpu_lts_20260413_b_RC00,build-id md589edbbe81c5b328a958fe628a9f2207d,781,691,048 字节,ELF x86-64 DYN,未 strip,clang/LLVM trunk)。该镜像携带完整的 C++ 符号表,因此下面的 pass 名称是 demangle 后的name()/getArgument()符号以及逐字的.rodata字符串,不是推断标签。其他 wheel 会不同。
摘要
本页是 libtpu 注册的每个内部编译器 pass 的名称索引:每个 pass 从 xla::HloPassInterface::name() 返回的逐字标识符、mlir::Pass 向 MLIR PassRegistry 注册的 -flag 参数,以及 xla::HloPassPipeline 在构造函数中携带的首字母大写阶段字符串。它是参考目录,不是算法页面:其价值在于按组去重的名称表,并交叉链接到记录每个 pass 主体的 Part-V 编译器页面。hlo-pass-registry.md 解释 pass 如何添加和排序,而本页回答:“我刚在 VLOG、xla_dump 目录或 --pass-pipeline= 规格中看到的这个字符串是什么,哪个页面记录了它?”
因为镜像未 strip,这些名称不是通过反汇编恢复,而是通过两种互补的符号枚举以接近 1:1 的保真度恢复。HLO 表面来自 _ZNK3xla*4nameEv 符号集:每个 XLA HLO pass 都会覆盖 HloPassInterface::name(),因此 demangle 这些符号即可直接枚举已注册的 HLO pass;其中 322 个是真实 pass,跨三个来源层级。MLIR 表面来自 create*Pass / Create*Pass 工厂符号以及 getArgument() pass-flag 字符串;它按方言清晰分组(StableHLO/CHLO/VHLO、Shardy、XTile、tpu、mosaic_sc、sparse_core、llo),底部以 242 个 tpu-* LLVM 后端 MachineFunction pass flag 收尾。少量首字母大写的 HloPassPipeline 阶段字符串和一个 phaseN_* 分离编译指标分类补全了清单。
读者会遇到的两种 name() 拼写按约定区分。HLO pass 的运行时名称是其类名的 kebab-case(ShardingPropagation → sharding-propagation);其中约 25 个已按字节确认为字面量字符串,其余在运行时构造,并根据类名推断(下方标记为 INFER)。MLIR pass 的名称是它交给 mlir::PassRegistration 的字面量;只要其 getArgument() 符号携带该字符串,就按字节确认。本页按流水线阶段组织名称:HLO-opt、sharding/SPMD、layout、fusion、collective、scheduling、MSA、host-offload、SparseCore,然后是 MLIR lowering 级联和 LLVM 后端 codegen 尾部;并为每组给出对应的文档 wiki 页面。状态/错误模板在 error-templates.md;建议提示字符串在 hint-strings.md;本页只负责 pass-name 索引和名称→页面交叉链接。
目录约定:
- 四类名称族以及每类如何恢复(
name()RTTI ·create*Pass工厂 ·getArgument()flag · pipeline stage 字符串)。 - **按阶段:**逐字注册名称、来源层级(
TTPU-private /Oopen-source XLA)和置信标签(按字节确认的字符串 vs 按约定推断)。 - 名称 → 文档页面交叉链接,覆盖每个有 Part-V 归属页面的分组。
- **真实缺口:**约 300 个 HLO
name()字符串按约定推断,逐 pass 的添加顺序归注册表页面所有,flag→pass gating 不在本页。
| 阶段组 | 族 | 数量 | 归属页面 |
|---|---|---|---|
| HLO-opt(expanders / canonicalizers / simplifiers) | HLO name() | ~95 | hlo-pre-passes.md, algebraic-simplifier.md |
| Sharding / SPMD / partition | HLO name() | ~16 | sharding-propagation.md, auto-sharding-spmd.md |
| Layout / tiling / mem-space | HLO name() | ~12 | layout-assignment.md |
| Fusion | HLO name() | ~13 | fusion-patterns.md, fusion-cost-model.md |
| Collective / async | HLO name() | ~40 | —(collective rewriting;未逐项分页) |
| Scheduling | HLO name() | ~14 | hlo-pass-registry.md |
| MSA(memory-space assignment) | symbol | ~8 | msa-overview.md |
| Host-offload | HLO name() | ~17 | — |
| While-loop / pipelining / misc | HLO name() | ~25 | loop-tiling-unrolling.md, optimization-barrier.md |
| SparseCore-offload(HLO-level) | HLO name() | ~10 | — |
| SparseCore tile-task(MLIR) | Create*Pass | 36 | lower-to-sparsecore-llvm.md |
| StableHLO / CHLO / VHLO legalizers | create*Pass/flag | ~40 | mhlo-xtile-tpu-lowering.md |
| Shardy(sdy)round-trip | create*Pass/flag | ~35 | sharding-propagation.md |
| XTile | create*Pass/flag | ~13 | mhlo-xtile-tpu-lowering.md |
tpu 方言(TensorCore vector IR) | create*Pass/flag | 14 | tpu-dialect-and-ops.md, tpu-to-llo-ods.md |
mosaic_sc 方言 | create*Pass | 3 | mosaic-overview.md |
llo / LLO bridge | create*Pass/flag | ~5 | lower-to-mlo-dma-bridge.md, llvmtpu-intrinsic-catalog.md |
| TPU LLVM-backend MachineFunction flags | cl::opt flag | 242 | —(codegen 尾部;仅 flags) |
| Pipeline-stage 字符串(monolithic) | HloPassPipeline ctor | ~15 | compile-phases.md |
| Separate-compilation phase metrics | .rodata token | 11 | compile-phases.md |
NOTE — 这些数量对符号族(HLO
name()、create*Pass、getArgument())而言是名称完整的,对 pipeline stage 而言是字符串确认的。任何 pipeline 内任意 pass 的添加顺序 不在本页:它归hlo-pass-registry.md和各阶段注册表所有。本页是名称索引;顺序索引是注册表。
名称如何恢复
四类名称族,四个提取表面。来源很重要,因为它决定置信度:demangle 后的符号对类是 CERTAIN,而它返回的运行时字符串只有在 .rodata 中也存在字面量时才是 CERTAIN。
HLO pass name ── _ZNK3xla*4nameEv symbol ── kebab(ClassName) ── string CERTAIN iff literal in .rodata, else INFER
MLIR pass flag ── create*Pass / Create*Pass factory symbol ── string CERTAIN where getArgument() carries it
MLIR pass flag ── PassIN…getArgumentEv symbol body ── byte-confirmed literal
pipeline stage ── HloPassPipeline ctor std::string at a .rodata offset ── byte-confirmed literal
phaseN_* metric ── .rodata token referenced from RegisterAllPhases ── byte-confirmed literal
```text
每张表使用的置信约定:
| 标签 | 含义 |
|---|---|
| `CERTAIN` | 已注册名称字符串是二进制中的字面量(`name()` `.rodata` 常量、`getArgument()` 字面量或 stage 字符串)——下方抽样确认。 |
| `HIGH` | pass *类*符号存在且已 demangle;运行时名称遵循 XLA 的 `CamelCase`→`kebab` 约定,但未定位到字面量(`name()` 主体在运行时构造)。 |
| `INFER` | 同 `HIGH`,用于大批 OSS 行,其中既未定位字面量,也未做逐 pass 反汇编;按约定可信,但没有字节锚点。 |
层级列:**`T`** = TPU-private(`xla::jellyfish::*`、`xla::tpu::sparse_core::*`、`xla::Tpu*`、`xla::megascale::*`)——Deepsea/Jellyfish 后端的私有表面;**`O`** = open-source XLA(`xla::*`),与 CPU/GPU 后端共享。
> **QUIRK —** HLO pass 的注册名称是*类名的 kebab-case,不是文件名或工厂名的 kebab-case*。`xla::jellyfish::TpuInt2AutoUpDownCaster` 注册为 `tpu-int2-auto-up-down-caster`:`Tpu` 前缀保留,命名空间不保留。若重新实现者从源文件名(`tpu_int2_caster.cc`)或命名空间路径派生名称,将与每个 `VLOG`/`xla_dump` 过滤器不匹配。
---
## HLO Pass 名称
HLO 表面是**322 个 pass**,跨三个来源层级:137 个 `xla::jellyfish::*`(TPU-private)、24 个 `xla::tpu::sparse_core::*`(SparseCore-HLO)、161 个 open-source `xla::*`,再加约 6 个 `megascale`/`sdy`/`spmd`/`conditional_opt`。下方按添加它们的 pipeline 阶段分组。名称是类从 `name()` 返回的 kebab-case;仅在定位到字面量时字符串为 `CERTAIN`,否则为 `INFER`。
### HLO-Opt — expanders / canonicalizers / simplifiers(Phase 1)
输入清理阶段:将高层 op 降低为原始算术,规范化形状和条件,折叠常量,并运行 algebraic simplifier。记录于 [`hlo-pre-passes.md`](../compiler/hlo-pre-passes.md);simplifier 本身记录于 [`algebraic-simplifier.md`](../compiler/algebraic-simplifier.md)。
| 注册名称 | 类 | 层级 |
|---|---|---|
| `tpu-int2-auto-up-down-caster` | `jellyfish::TpuInt2AutoUpDownCaster` | T |
| `algebraic-simplifier` / TPU superset | `jellyfish::TpuAlgebraicSimplifier`, `AlgebraicSimplifier` | T/O |
| `batch-norm-expander` | `BatchNormExpander` | O |
| `cholesky-expander` | `TpuCholeskyExpander`, `CholeskyExpander` | T/O |
| `qr-expander` / `eigh-expander` | `TpuQrExpander`/`TpuEighExpander` (+ OSS) | T/O |
| `triangular-solve-expander` | `TpuTriangularSolveExpander`, `TriangularSolveExpander` | T/O |
| `fft-expander`, `lu-decomposition-expander` | `FftExpander`, `LuDecompositionExpander` | O |
| `dynamic-padder`, `dynamic-index-splitter`, `dynamic-dimension-simplifier` | `DynamicPadder` etc. | O |
| `scatter-expander` / `gather-expander` | `TpuScatterExpander`/`TpuGatherExpander` (+ OSS) | T/O |
| `comparison-expander`, `logistic-expander`, `optimization-barrier-expander` | `ComparisonExpander` etc. | O |
| `dot-decomposer`, `dot-canonicalizer`, `dot-algorithm-normalizer` | `DotDecomposer`, `jellyfish::Dot*` | O/T |
| `x-precision-rewriter` | `jellyfish::XPrecisionRewriter` (x6/x9/x128 dot) | T |
| `tpu-hlo-precision-tracer`, `float-type-correction` | `jellyfish::TpuHloPrecisionTracer`, `FloatTypeCorrection` | T |
| `tpu-broadcast-rewriter`, `tpu-degenerate-dimension-rewriter` | `jellyfish::TpuBroadcastRewriter` etc. | T |
| `tpu-reduce-rewriter`, `tpu-reduce-window-{rewriter,normalizer}` | `jellyfish::TpuReduce*` | T |
| `tpu-convolution-rewriter`, `convolution-folding` | `jellyfish::TpuConvolutionRewriter`, `ConvolutionFolding` | T |
| `tpu-call-inliner` (`-must-fuse`/`-inner-must-fuse`/`-non-must-fuse`) | `jellyfish::TpuCallInliner` | T |
| `tpu-hlo-support-checker` | `TpuHloSupportChecker`(规范验收测试) | T |
| `hlo-dce` / `hlo-module-dce` / `hlo-cse` | `HloDCE`/`HloModuleDCE`/`HloCSE`(阶段间重跑) | O |
| `flatten-call-graph`, `call-inliner`, `map-inliner` | `FlattenCallGraph` etc. | O |
| `float-normalization`, `bfloat16-propagation` | `FloatNormalization`, `BFloat16Propagation` | O |
| `gmm`, `ragged-dot-expander`, `flash-attention` | `jellyfish::Gmm`/`RaggedDotExpander`/`FlashAttention` | T |
> **NOTE —** `jellyfish::Tpu*` rewriter 是相同根名称 OSS pass 的*超集*,不是替代品;`TpuAlgebraicSimplifier` 会重跑 OSS `AlgebraicSimplifier` 规则集并加上 TPU-only 规则。两个类都有不同的 `name()` 符号,因此 `VLOG` 过滤器会看到两个字符串。
### Sharding / SPMD / Partition(Phase 2-3)
Sharding propagation 和 SPMD partitioning。`sharding-propagation` 已按字节确认(字面量出现三次,分别对应 propagation、dump 和 verifier 路径中的注册)。记录于 [`sharding-propagation.md`](../compiler/sharding-propagation.md) 和 [`auto-sharding-spmd.md`](../compiler/auto-sharding-spmd.md)。
| 注册名称 | 类 | 层级 |
|---|---|---|
| `sharding-propagation` | `ShardingPropagation`(×3 registration) | O |
| `auto-sharding` | `AutoSharding`, `TpuAutoSharding` | O/T |
| `tpu-spmd-partitioning` | `jellyfish::TpuSpmdPartitioner` | T |
| `tpu-partition-assignment` | `jellyfish::TpuPartitionAssignment` | T |
| `tpu-spmd-concat-rewriter` | `jellyfish::TpuSpmdConcatRewriter` | T |
| `tpu-sparse-core-hierarchical-spmd-partitioning` | `sparse_core::SparseCoreHierarchicalSpmdPartitioner` | T |
| `whole-graph-manual-pass` | `WholeGraphManualPass` | O |
| `spmd-prepare`, `spmd-partitioner` | `spmd::SpmdPrepare`/`SpmdPartitioner` | O |
| `shardy-xla` | `sdy::ShardyXLA`(Shardy frontend import) | O |
| `hlo-domain-isolator` / `hlo-domain-remover` | sharding-domain brackets | O |
| `canonicalize-all-gather-for-cse` | `CanonicalizeAllGatherForCSE` | O |
### Layout / Tiling / Memory-Space(Phase 4-5)
Layout assignment、TPU tiling propagation 和 memory-space propagation。记录于 [`layout-assignment.md`](../compiler/layout-assignment.md)。
| 注册名称 | 类 | 层级 |
|---|---|---|
| `layout-assignment` | `LayoutAssignment` | O |
| `tpu-tiling-propagation`, `tpu-tiling-rewriter` | `jellyfish::TpuTilingAssignment`/`TpuTilingRewriter` | T |
| `tpu-peephole-optimize-layouts` | `jellyfish::TpuPeepholeOptimizeLayouts` | T |
| `window-config-assignment` (`-via-service`) | `jellyfish::WindowConfigAssignment[ViaService]` | T |
| `memory-space-propagation` | `MemorySpacePropagation` | O |
| `copy-insertion`, `tpu-copy-normalizer` | `CopyInsertion`, `jellyfish::TpuCopyNormalizer` | O/T |
| `convert-memory-placement-to-internal-annotations` | `ConvertMemoryPlacementToInternalAnnotations` | O |
| `optimize-input-output-buffer-alias` | `OptimizeInputOutputBufferAlias` | O |
### Fusion(Phases "Pre main fusion" / "Main fusion" / "Post main fusion")
Fusion 桶。`tpu-instruction-fusion`(主 pass,使用 `TpuPriorityFusionQueue`)在运行时构造其名称:类符号 `xla::jellyfish::TpuInstructionFusion::name()` 位于 `0x12e037a0`,但不携带 `.rodata` 字面量,因此该名称按约定为 `HIGH` 置信。记录于 [`fusion-patterns.md`](../compiler/fusion-patterns.md) 和 [`fusion-cost-model.md`](../compiler/fusion-cost-model.md)。
| 注册名称 | 类 | 层级 |
|---|---|---|
| `tpu-instruction-fusion` | `jellyfish::TpuInstructionFusion` (`name()` @ `0x12e037a0`) | T |
| `tpu-multi-output-fusion` | `jellyfish::TpuMultiOutputFusion` | T |
| `instruction-fusion` / `multi-output-fusion` / `megacore-fusion` | OSS `InstructionFusion` etc. | O |
| `tpu-loop-fusion-enhancer` | `jellyfish::TpuLoopFusionEnhancer` | T |
| `mosaic-fusion` | `jellyfish::MosaicFusion`(Pallas/Mosaic,`HloPassFix`) | T |
| `tpu-subgraph-simplifier` | `jellyfish::PostFusionTpuSubgraphSimplifier` | T |
| `user-guided-fusion-id-assigner` | `jellyfish::UserGuidedFusionIdAssigner` | T |
| `wrap-fusion-output-for-debug` | `jellyfish::WrapFusionOutputForDebug` | T |
| `fusion-tuple-deduplicator`, `masked-fusion-iteration-skipper` | `jellyfish::FusionTupleDeduplicator` etc. | T |
| `tpu-trivial-fusion-remover` / `tpu-trivial-instruction-unfuser` | `jellyfish::TpuTrivial*` | T |
### Collective / Async Rewriting(Phase 4 起)
All-reduce / all-gather / reduce-scatter / all-to-all combiner、legalizer 和 async-collective creator,包括 MegaScale cross-slice 家族。未逐项分页;本节是名称索引。
| 注册名称 | 类 | 层级 |
|---|---|---|
| `tpu-all-reduce-legalizer`, `tpu-all-reduce-scatter-fusion` | `jellyfish::TpuAllReduceLegalizer` etc. | T |
| `tpu-all-gather-combiner`, `tpu-all-gather-normalizer` | `jellyfish::TpuAllGather*` | T |
| `all-reduce-combiner` / `all-reduce-reassociate` / `all-reduce-simplifier` | OSS `AllReduce*` | O |
| `all-gather-combiner` / `all-gather-decomposer` / `all-gather-simplifier` | OSS `AllGather*` | O |
| `tpu-async-collective-creator` | `jellyfish::TpuAsyncCollectiveCreator` | T |
| `async-collective-merger` / `async-collective-fusion` / `async-op-thread-renamer` | `jellyfish::Async*` | T |
| `collective-pipeliner`, `collective-permute-combiner`/`-cse`/`-motion` | OSS `Collective*` | O |
| `reduce-scatter-decomposer` / `reduce-scatter-legalizer` | OSS + `jellyfish::ReduceScatterLegalizer` | O/T |
| `ragged-all-to-all-expander`, `tpu-all-to-all-compaction` | `jellyfish::Ragged*`/`TpuAllToAllCompaction` | T |
| `tpu-quantized-all-reduce-backend-config-setter` | `jellyfish::TpuQuantizedAllReduceBackendConfigSetter` | T |
| `barrier-assignment`, `async-barrier-creator` | `jellyfish::BarrierAssignment` etc. | T |
| MegaScale: `tpu-all-reduce-merger`, `cross-slice-legalizer`/`-rewrites` | `megascale::compiler::*` | T |
### Scheduling(Phase 7: "final_scheduler" / "async_scheduling")
Latency-hiding scheduler 家族。三个 scheduler `name()` 字符串已按字节确认为带短横线的字面量:`latency-hiding-layer-scheduler`、`sparsecore-latency-hiding-scheduler` 和 `legalize-scheduling-annotations`(各命中一次)。OSS `latency-hiding-scheduler` 名称在运行时构造:`.rodata` 中只存在两个 TPU 前缀形式里的子串 `latency-hiding-scheduler`,因此该行为 `HIGH`,不是 `CERTAIN`。Scheduler 的放置归 [`hlo-pass-registry.md`](../compiler/hlo-pass-registry.md) 所有。
| 注册名称 | 类 | 层级 |
|---|---|---|
| `latency-hiding-scheduler` | `LatencyHidingScheduler` | O |
| `latency-hiding-layer-scheduler` | `jellyfish::LatencyHidingLayerScheduler` (`name()` @ `0x10b85160`) | T |
| `sparsecore-latency-hiding-scheduler` | `sparse_core::SparseCoreLatencyHidingScheduler` (`name()` @ `0x13077ea0`) | T |
| `legalize-scheduling-annotations` | `LegalizeSchedulingAnnotations` (`name()` @ `0x12e9a060`) | O |
| `check-no-data-dependency-in-scheduling-annotations` | `CheckNoDataDependencyInSchedulingAnnotations` | O |
| `annotate-traceable-loops` | `jellyfish::TpuAnnotateTraceableLoops` | T |
| `while-loop-all-reduce-code-motion` | `WhileLoopAllReduceCodeMotion` | O |
| `hlo-memory-scheduler`, `hlo-memory-scheduler-with-brkga-fallback`, `hlo-descheduler` | OSS schedulers | O |
| `async-op-scheduler`, `barna-core-feed-scheduler`, `slack-in-trough`, `constant-deferring` | `jellyfish::*` / OSS | T/O |
> **QUIRK —** `final_scheduler` 和 `async_scheduling` 是 **pipeline *stage* 名称**,不是 pass 名称:它们是两个 scheduler pipeline 的 `HloPassPipeline` 构造函数字符串,位于下方的 [pipeline-stage 小节](#pipeline-stage-names)。这些 pipeline *内部*的 pass 是上面的 `*-scheduler` 行。不要把 pipeline 标签和其中包含的 pass 混淆。
### MSA — Memory-Space Assignment(调度之后)
MSA 是一个引擎,不是带有 `name()` 的 `HloPass`;它由 `xla::jellyfish::RunMemorySpaceAssignment` 驱动。记录于 [`msa-overview.md`](../compiler/msa-overview.md)。
| 符号 / 名称 | 层级 |
|---|---|
| `memory_space_assignment::MemorySpaceAssignment`(引擎) | O |
| `MemorySpaceAssignmentBestFitRepacker`(telamalloc-class repacker) | O |
| `jellyfish::RunMemorySpaceAssignment`(driver fn) | T |
| `jellyfish::ScopedMemoryAssignment`(`scoped-vmem` anchor) | T |
| `sparse_core::SparseCoreMemorySpaceAssignment` | T |
| `tpu-memory-precolor`(`jellyfish::TpuMemoryPrecolor`) | T |
| `jellyfish::PlumbVmemInfo`, `PinPrecoloring`, `PinUserAnnotatedInstructions` | T |
### Host-Offload(Phase 5-6)
Host-offload 分析和 asyncification。未逐项分页。
| 注册名称 | 类 | 层级 |
|---|---|---|
| `host-offloader` / `host-offload-legalize` / `host-offloading-prepare` | OSS `HostOffload*` | O |
| `host-offloading-layout-analysis`, `annotate-host-compute-offload` | OSS | O |
| `host-compute-asyncifier`, `host-memory-transfer-asyncifier` | OSS | O |
| `add-random-host-offloading` | `jellyfish::AddRandomHostOffloading`(debug) | T |
| `host-offloading-lookahead`, `host-offload-normalization` | `jellyfish::*` | T |
| `host-collective-rewriter`, `host-transfer-channel-id-assigner` | `jellyfish::*` | T |
| `infeed-decomposer` / `infeed-flattener` / `outfeed-decomposer` | `jellyfish::*` | T |
| `infeed-token-propagation` | `InfeedTokenPropagation` | O |
### While-Loop / Pipelining / Misc
While-loop 简化、展开、rematerialization,以及其他 TPU rewrite。Loop transform 记录于 [`loop-tiling-unrolling.md`](../compiler/loop-tiling-unrolling.md);opt-barrier 处理记录于 [`optimization-barrier.md`](../compiler/optimization-barrier.md)。
| 注册名称 | 类 | 层级 |
|---|---|---|
| `while-loop-simplifier` / `while-loop-constant-sinking` / `while-loop-invariant-code-motion` | OSS `WhileLoop*` | O |
| `while-loop-unroller` / `while-loop-pipeline-unroller` / `while-loop-trip-count-annotator` | OSS | O |
| `while-loop-concat-code-motion` / `-expensive-invariant-code-motion` / `-fusible-sinking` | OSS | O |
| `hlo-rematerialization` / `hlo-rematerialization-with-stats` | OSS | O |
| `add-original-value` | `AddOriginalValue`(provenance metadata) | O |
| `tpu-initial-weight-broadcaster` | `jellyfish::TpuInitialWeightBroadcaster` | T |
| `tpu-indexed-instruction-analyzer` | `jellyfish::TpuIndexedInstructionAnalyzer` | T |
| `convert-frontend-attributes-to-backend-config` | `jellyfish::ConvertFrontendAttributesToBackendConfig`(HLO domain 中最后一个) | T |
| `cross-slice-code-motion`, `pred-packing`, `operand-span-exposure` | `jellyfish::*` | T |
| `always-crash` | `jellyfish::AlwaysCrash`(testing,由 `--xla_tpu_always_crash` gate) | T |
| `mock-while-loop-bounding` | `jellyfish::MockWhileLoopBounding`(testing) | T |
| `conditional-code-motion` | `conditional_opt::ConditionalCodeMotion` | O |
### SparseCore-Offload(HLO-level)
在 SparseCore MLIR lowering *之前*运行的 HLO-level SparseCore offloader。24 个 `xla::tpu::sparse_core::*` HLO pass(embedding / sparse-dense matmul decomposer)属于 SparseCore-HLO 层级;MLIR tile-task pass 位于 [SparseCore MLIR 小节](#sparsecore-tile-task-passes)。记录于 [`lower-to-sparsecore-llvm.md`](../compiler/lower-to-sparsecore-llvm.md)。
| 注册名称 | 类 | 层级 |
|---|---|---|
| `sparse-core-compute-offloader` | `jellyfish::SparseCoreComputeOffloader` | T |
| `sparse-core-collective-offload` (+ `-remover`/`-aggregator`) | `jellyfish::SparseCoreCollective*` | T |
| `offload-gather-to-sparse-core` / `offload-scatter-to-sparse-core` / `offload-sort-to-sparse-core` | `jellyfish::Offload*ToSparseCore` | T |
| `sparse-core-queue-assignment` / `sparse-core-queue-legalizer` | `jellyfish::SparseCoreQueue*` | T |
| `sparse-core-mismatch-detector`, `aggregate-data-dependent-sc-ops` | `jellyfish::*` | T |
| 24 `sparse_core::*` HLO passes(embedding/SDM decomposers) | `tpu::sparse_core::*` | T |
---
## MLIR Pass 名称
MLIR 表面从 `create*Pass`/`Create*Pass` 工厂符号和 `getArgument()` pass-flag 字面量恢复。它按方言拆分,并形成 4 阶段 lowering 级联。下方名称是 `getArgument()` 字面量(在 `getArgument()` 符号携带它的位置按字节确认)。
```text
StableHLO ── createStablehloLowerToXtilePass ──> XTile
XTile ── vector-layout / tiling passes ──> tpu (TensorCore vector IR)
tpu ── createLowerToLLOPass ──> LLO ── tpu-bundle-packer / tpu-encode-mcinst-bundles ──> ISA bundles
tpu ── (SparseCore fork) sc 36 passes ──> Mlo ── createLowerToSparseCoreLlvmPass ──> LLVM-TPU IRStableHLO / CHLO / VHLO Legalizers
输入 wire-format 合法化。记录于 mhlo-xtile-tpu-lowering.md。
| Pass-arg flag | 工厂 |
|---|---|
chlo-legalize-to-stablehlo / -hlo / -high-level-mhlo | createChloLegalize*Pass |
stablehlo-canonicalize-dynamism | createStablehloCanonicalizeDynamismPass |
stablehlo-complex-math-expander | createStablehloComplexMathExpanderPass |
stablehlo-convert-to-signless | createStablehloConvertToSignlessPass |
stablehlo-legalize-quant-to-math | createStablehloLegalizeQuantToMathPass |
stablehlo-legalize-to-{hlo,linalg,vhlo} | createStablehloLegalizeTo*Pass |
stablehlo-refine-{arguments,shapes} | createStablehloRefine*Pass |
vhlo-to-version, vhlo-legalize-to-stablehlo | createVhlo*Pass |
legalize-tf, legalize-tf-{collective,communication} | createLegalizeTFPass |
stablehlo-ext-*(8 个变体:canonicalize-dynamism、refine-shapes、sink-constants-to-control-flow、…) | extension passes |
Shardy(sdy)Round-Trip Sharding
Shardy import/export round-trip。记录于 sharding-propagation.md。
| Pass-arg flag | 工厂 |
|---|---|
sdy-import-shardings / -import-constants / -import-func-calls | xla::sdy::createImport*Pass |
sdy-export-ops / -export-named-computations | createExportOpsPass etc. |
sdy-round-trip-{dedup-meshes,export-ops} | createSdyRoundTrip*Pass |
sdy-apply-sharding-constraints, -close-shardings, -collective | sdy passes |
sdy-insert-explicit-reshards, -reshard, -reshard-to-collectives | sdy passes |
sdy-flatten-call-graph, -inline-meshes, -lift-inlined-meshes | sdy passes |
总计约 35 个 sdy-* flags(constant-or-scalar-{merger,splitter}, remove-*, manual-axes-cleanup, …) | sdy passes |
XTile
StableHLO → XTile lowering 层,以及 XTile-CPU bufferization/vectorization pass。记录于 mhlo-xtile-tpu-lowering.md。
| Pass-arg flag | 工厂 |
|---|---|
stablehlo-lower-to-xtile | xla::xtile::createStablehloLowerToXtilePass |
stablehlo-lower-to-arith | createStablehloLowerToArithPass |
xtile-verify-legal-ops | createVerifyLegalXTileOpsPass |
xtile-cpu-*(bufferization、fuse-elementwise、linalg-elementwise-to-vector、lower-xtile-entry、shlo-to-vector、vector-to-scalar、unpack-sub-byte-vector-write、…) | XTile-CPU passes |
tpu 方言 — TensorCore Vector IR(14 个 create*Pass)
TensorCore vector-layout 和 lowering pass。lower-to-llo 和 lower-to-mlo 已按字节确认(已定位 getArgument() 字面量)。记录于 tpu-dialect-and-ops.md 和 tpu-to-llo-ods.md。
| Pass-arg flag | 工厂 |
|---|---|
lower-to-llo | mlir::tpu::createLowerToLLOPass (getArgument() @ 0x112041c0) |
lower-to-mlo | mlir::tpu::createLowerToMloPass |
tpu-tiling-propagation / tiling-propagation | createTilingPropagationPass (getArgument() @ 0x132e0d80) |
tpu-apply-vector-layout | createApplyVectorLayoutPass |
tpu-infer-vector-layout / infer-vector-layout | createInferVectorLayoutPass |
tpu-infer-memref-layout | createInferMemRefLayoutPass |
tpu-relayout-insertion / relayout-insertion | createRelayoutInsertionPass |
tpu-canonicalize-{memory-space,mosaic,operations} | createCanonicalize*Pass |
(无 flag)createConvertIntegerMemrefsPass, createDebugAssertInsertionPass, createLogicalToPhysicalDeviceIdPass, createPreCanonicalizationOptimizationPass | tpu passes |
mosaic_sc 方言 — Mosaic-for-SparseCore(3 个 create*Pass)
记录于 mosaic-overview.md。
| 工厂 |
|---|
mlir::mosaic_sc::createApplyVectorLayoutPass |
mlir::mosaic_sc::createInferVectorLayoutPass |
mlir::mosaic_sc::createInsertRelayoutPass |
SparseCore Tile-Task Passes(36 个 Create*Pass)
xla::tpu::sparse_core::Create*Pass 家族:SparseCore tile-task lowering 链(bounds-check、tile outlining、allocation、prefetch scheduling、LLVM lowering)。记录于 lower-to-sparsecore-llvm.md。按功能命名,而不是平铺 dump:
| 子组 | 代表性 Create*Pass 工厂 |
|---|---|
| Bounds / safety checks | AddStreamBoundChecks, IndirectStreamBoundsCheck, LoadStoreBoundsCheck, InsertRunIdCheck, LlvmIntToPtrSafety, SyncFlagVerifier |
| Tile outlining / structure | TileTaskOutlining, SinkOpsIntoTileTask, ExpandTiledMemRefs, TileTaskArgumentsSpill, InsertClearTileIbuf |
| Allocation | AllocationAssignment, GlobalAllocationAssignment, TileOverlayAllocation, LegalizeHbmGlobals, PrepareHbmSpill |
| Scheduling / prefetch | ScheduleTileTaskPrefetch, LowerPrefetchTileTask, DelaysInsertion, TileBarrierInsertion |
| Lowering | LowerToSparseCoreLlvm, LowerDynamicDimensionSize, LowerSequencerFunctions, PackedOperandsLowering, AlgebraicDecomposer |
| Misc | HardwareTracing, SflagWaitInstrumentation, EmitTileOverlayHandler, ScsOverlaysTrampoline, MemrefAliasScopeAnnotation, PackUnpackCleaner, VectorUnroll, LoopUnroll, SparseCoreCanonicalizer, TileHostTransfer, LegalizeMegaChipModule |
| Pass-arg flags | legalize-hbm-globals, legalize-megachip-module, lower-to-mlo, lower-to-sc, lower-to-sparse-core-llvm |
llo / LLO Bridge
在 bundle packing 之前关闭 MLIR 级联的 LLO-extension elimination pass。记录于 lower-to-mlo-dma-bridge.md 和 llvmtpu-intrinsic-catalog.md。
| 工厂 / 类 |
|---|
mlir::llo::createEliminateLLOExtensionsPass / EliminateLLOExtensionsPass |
TPU LLVM-Backend MachineFunction Pass Flags(242)
Codegen 栈底部,位于 LLO 之下:TPU LLVM target 的 machine-level pass,暴露为 242 个 tpu-* cl::opt flag 字符串。只恢复了 flag 字符串;它们背后的 llvm::Pass 子类符号不在本页(codegen 尾部)。按功能分组,而不是平铺 dump:
| 功能 | 代表性 flags |
|---|---|
| Scheduling | tpu-machine-scheduler[-fast], tpu-misched[-dcpl], tpu-fifo-scheduler, tpu-use-{fifo-sched,swing-modulo-sched}, tpu-pipeliner[-analysis/-info/-strategy/-verbose], tpu-pipeline-dovetailing, tpu-slack-lazy-backtrack |
| Super-passes | tpu-ssa-super-pass, tpu-enable-pipeliner-super-pass, tpu-enable-overlayer-passes |
| Bundle packing / encoding | tpu-bundle-packer[-maintain-order/-maintain-topologial-order], tpu-encode-mcinst-bundles, tpu-post-bundle-lower, tpu-print-{mcinst,sc-mcinst}-encodings |
| Register alloc / spill | tpu-register-prepare, tpu-round-robin-registers, tpu-pre-spill[-per-fill], tpu-spill-to-dreg, tpu-opt-spill-to-dreg, tpu-crosscall-spill-packer, tpu-fixed-{vregs,maskregs}, tpu-liveness-*, tpu-remat-pass |
| ISel / lowering | tpu-isel, tpu-early-post-isel, tpu-opt-prepare, tpu-codegen-prep[are], tpu-frame-lowering, tpu-gep-lowering, tpu-rotate-lowering, tpu-memop-intrinsics, tpu-vliw-prepare, tpu-mem-alloc |
| Verifiers | tpu-verifier, tpu-machine-memoperands-verifier, tpu-validate-{inputs,session-inputs,function-sizes}, tpu-fatal-{function-size-assertions,mem-alloc-error,verifier-error} |
| V1-island / TF-bridge | tpu-v1-island-{coarsening,inlining,outlining}, tpu-cluster-formation, tpu-device-propagation, tpu-merge-variables-with-execute, tpu-colocate-{composite-resource-ops,splits} |
GOTCHA — 若干
tpu-*字符串是双表面的:tpu-call-inliner、tpu-instruction-fusion、tpu-int2-auto-up-down-caster、tpu-tiling-propagation、tpu-partition-assignment既作为 HLOname()字符串(上方)出现,也作为这里的tpu-*flag 参数出现。它们是同一个逻辑 pass 暴露在两个注册层级上。把tpu-*flag 和 HLOname()字符串算作互不相交集合会重复计数;应将 HLO-level pass 视为规范名称,flag 视为它的 CLI alias。
Pipeline-Stage 名称 {#pipeline-stage-names}
两套阶段分类并存。Monolithic compile 用首字母大写字符串(传给 pipeline 构造函数的 std::string)标记其 HloPassPipeline 阶段;separate compile(tpu_phase_functions.cc)用 phaseN_* 指标 token 标记其阶段。它们不能互换。两者都记录于 compile-phases.md。
Monolithic HloPassPipeline Stage Strings
| Stage string | File-off | 角色 |
|---|---|---|
pre-optimization | 0x85ea920 | PreOptimizationPipeline(Phase 1) |
post-optimization | 0x85ea85e | PostOptimizationPipeline |
Pre layout assignment | 0x84ef7bf | pre-layout HLO |
Layout assignment / XLA::JF Layout Assignment | 0x84ef7d5 / 0x84efbdf | LayoutAssignment |
Post layout assignment | 0x84ef6db | post-layout HLO |
HLO passes through layout assignment | 0x84ef757 | HloOptimizeThroughLayoutAssignment |
HLO passes after layout assignment | 0x84ef6f2 | HloOptimizeAfterLayoutAssignment |
Pre main fusion | 0x10963b16 | PostLayoutPreFusionHloOptimize |
Main fusion / XLA::TPU main fusion pass | — / 0x8522e3e | TpuInstructionFusion bucket |
Post main fusion / JF after main fusion | 0x85f7605 / 0x85f7616 | PostMainFusionHloOptimize |
final_scheduler | (SSO strcpy) | 基础 memory-schedule pipeline(Phase 7) |
async_scheduling | 0x865dcab | LHS + async-overlap pipeline(Phase 7) |
final_scheduled | 0x8713d6a | post-scheduling marker |
simplification after layout assignment | 0x84ef715 | post-layout simplify |
Separate-Compilation Phase Metric Tokens
Multi-phase compile 用 phaseN_* token 命名其阶段(separate-compilation timing 表中的 metric key)。phase1_hlo_opts 和 phase3_linking 已按字节确认为从 xla::TpuCompiler::RegisterAllPhases @ 0xf849ec0 引用的字面量。“TLP” = TPU-Level-Program。
| Token | 角色 |
|---|---|
phase0_stablehlo_to_hlo | StableHLO → HLO import |
phase1_hlo_opts | HLO optimization(上方 HLO passes) |
phase2a_tlp_lowering | HLO → MLIR-tpu(TPU-Level-Program)lowering |
phase2b_deduped_lowering | dedup'd per-chunk lowering |
phase3_linking (+ _test_only) | HLO-module linking |
phase2_compilation_{status,time} | phase-2 status/timing metric fields |
phase0_{transfer_dim,shard_size_bytes}, phase1_transfer_dim | per-phase transfer/shard metric fields |
NOTE —
phaseN_*token 是指标 key,不是 pass 名称:它们出现在 timing/transfer 表中,绝不会作为读者可以运行的 pass 出现。VLOG或xla_dump过滤器会匹配上方的 stage 字符串和前面的 pass 名称;除非读取 separate-compilation timing 路径,否则不会匹配phaseN_*。锚定phase1_hlo_opts/phase3_linking的RegisterAllPhases符号是注册入口,不是逐阶段 pass 主体。
未确认 / 不在本页
诚实列出范围限制,便于重新实现者知道需要重新推导什么:
- 约 300 个 HLO
name()字符串按约定推断(INFER),未按字节确认。name()主体在运行时构造字符串(static const 或 string-builder);只有约 25 个带短横线名称作为.rodata字面量存在(CERTAIN行)。类符号是确定的;精确返回字符串遵循CamelCase→kebab约定,但未逐 pass 定位。 - 每个 pass 的添加顺序不在这里。 本页是名称索引;顺序归
hlo-pass-registry.md(Phase-1..6 序列)、scheduler placement 和 fusion-bucket order 所有。 - flag → pass gating 不在这里。 哪个
xla_tpu_*flag 启用哪个AddPass<T>是RunHloPasses(0x1093a420)中的.text分支,不是本页恢复的字符串。 - 242 个
tpu-*flag 的llvm::Pass类符号未恢复:只恢复了cl::optflag 字符串。 - per-TpuVersion(jxc/pxc/vxc/glc/gfc)pass-set delta 未隔离:HAL 模块按版本命名,但 v3 与 v6e 哪些 pass 被跳过是
Target::tpu_version()分支,本页未追踪。
交叉引用
hlo-pass-registry.md— pass-add 注册表;本页是它的名称索引,该页是它的顺序索引compile-phases.md— stage 字符串所标记的 monolithic-vs-separate 阶段分类algebraic-simplifier.md—algebraic-simplifier/TpuAlgebraicSimplifier行的主体sharding-propagation.md—sharding-propagation(按字节确认 ×3)和 sdy round-trip passes 的主体auto-sharding-spmd.md—auto-sharding/tpu-spmd-partitioning行的主体layout-assignment.md—layout-assignment/tpu-tiling-propagation行的主体fusion-patterns.md—tpu-instruction-fusion和 fusion bucket 的主体loop-tiling-unrolling.md— while-loop unroll / pipeline 行的主体optimization-barrier.md—optimization-barrier-expander行的主体msa-overview.md— memory-space 行背后的 MSA 引擎mhlo-xtile-tpu-lowering.md— StableHLO/XTile/tpuMLIR lowering 级联lower-to-sparsecore-llvm.md— 36 个 SparseCoreCreate*Passlowering 链tpu-to-llo-ods.md—lower-to-llo/lower-to-mlo主体error-templates.md— error/status template 目录(对比:本页负责 pass 名称,不负责错误文本)hint-strings.md— advisory-hint 目录(对比:本页负责 pass 名称,不负责补救方案)overview.md— Part XI runtime section map