Skip to content

内部 Pass 名称目录

本页中的所有地址和偏移都适用于 libtpu-0.0.40-cp314 wheel 中的 libtpu.solibtpu_lts_20260413_b_RC00,build-id md5 89edbbe81c5b328a958fe628a9f2207d,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 如何添加和排序,而本页回答:“我刚在 VLOGxla_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、tpumosaic_scsparse_corello),底部以 242 个 tpu-* LLVM 后端 MachineFunction pass flag 收尾。少量首字母大写的 HloPassPipeline 阶段字符串和一个 phaseN_* 分离编译指标分类补全了清单。

读者会遇到的两种 name() 拼写按约定区分。HLO pass 的运行时名称是其类名的 kebab-caseShardingPropagationsharding-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 字符串)。
  • **按阶段:**逐字注册名称、来源层级(T TPU-private / O open-source XLA)和置信标签(按字节确认的字符串 vs 按约定推断)。
  • 名称 → 文档页面交叉链接,覆盖每个有 Part-V 归属页面的分组。
  • **真实缺口:**约 300 个 HLO name() 字符串按约定推断,逐 pass 的添加顺序归注册表页面所有,flag→pass gating 不在本页。
阶段组数量归属页面
HLO-opt(expanders / canonicalizers / simplifiers)HLO name()~95hlo-pre-passes.md, algebraic-simplifier.md
Sharding / SPMD / partitionHLO name()~16sharding-propagation.md, auto-sharding-spmd.md
Layout / tiling / mem-spaceHLO name()~12layout-assignment.md
FusionHLO name()~13fusion-patterns.md, fusion-cost-model.md
Collective / asyncHLO name()~40—(collective rewriting;未逐项分页)
SchedulingHLO name()~14hlo-pass-registry.md
MSA(memory-space assignment)symbol~8msa-overview.md
Host-offloadHLO name()~17
While-loop / pipelining / miscHLO name()~25loop-tiling-unrolling.md, optimization-barrier.md
SparseCore-offload(HLO-level)HLO name()~10
SparseCore tile-task(MLIR)Create*Pass36lower-to-sparsecore-llvm.md
StableHLO / CHLO / VHLO legalizerscreate*Pass/flag~40mhlo-xtile-tpu-lowering.md
Shardy(sdy)round-tripcreate*Pass/flag~35sharding-propagation.md
XTilecreate*Pass/flag~13mhlo-xtile-tpu-lowering.md
tpu 方言(TensorCore vector IR)create*Pass/flag14tpu-dialect-and-ops.md, tpu-to-llo-ods.md
mosaic_sc 方言create*Pass3mosaic-overview.md
llo / LLO bridgecreate*Pass/flag~5lower-to-mlo-dma-bridge.md, llvmtpu-intrinsic-catalog.md
TPU LLVM-backend MachineFunction flagscl::opt flag242—(codegen 尾部;仅 flags)
Pipeline-stage 字符串(monolithic)HloPassPipeline ctor~15compile-phases.md
Separate-compilation phase metrics.rodata token11compile-phases.md

NOTE — 这些数量对符号族(HLO name()create*PassgetArgument())而言是名称完整的,对 pipeline stage 而言是字符串确认的。任何 pipeline 内任意 pass 的添加顺序 在本页:它归 hlo-pass-registry.md 和各阶段注册表所有。本页是名称索引;顺序索引是注册表。


名称如何恢复

四类名称族,四个提取表面。来源很重要,因为它决定置信度:demangle 后的符号对CERTAIN,而它返回的运行时字符串只有在 .rodata 中也存在字面量时才是 CERTAIN

text
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 IR

StableHLO / CHLO / VHLO Legalizers

输入 wire-format 合法化。记录于 mhlo-xtile-tpu-lowering.md

Pass-arg flag工厂
chlo-legalize-to-stablehlo / -hlo / -high-level-mhlocreateChloLegalize*Pass
stablehlo-canonicalize-dynamismcreateStablehloCanonicalizeDynamismPass
stablehlo-complex-math-expandercreateStablehloComplexMathExpanderPass
stablehlo-convert-to-signlesscreateStablehloConvertToSignlessPass
stablehlo-legalize-quant-to-mathcreateStablehloLegalizeQuantToMathPass
stablehlo-legalize-to-{hlo,linalg,vhlo}createStablehloLegalizeTo*Pass
stablehlo-refine-{arguments,shapes}createStablehloRefine*Pass
vhlo-to-version, vhlo-legalize-to-stablehlocreateVhlo*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-callsxla::sdy::createImport*Pass
sdy-export-ops / -export-named-computationscreateExportOpsPass etc.
sdy-round-trip-{dedup-meshes,export-ops}createSdyRoundTrip*Pass
sdy-apply-sharding-constraints, -close-shardings, -collectivesdy passes
sdy-insert-explicit-reshards, -reshard, -reshard-to-collectivessdy passes
sdy-flatten-call-graph, -inline-meshes, -lift-inlined-meshessdy 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-xtilexla::xtile::createStablehloLowerToXtilePass
stablehlo-lower-to-arithcreateStablehloLowerToArithPass
xtile-verify-legal-opscreateVerifyLegalXTileOpsPass
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-llolower-to-mlo 已按字节确认(已定位 getArgument() 字面量)。记录于 tpu-dialect-and-ops.mdtpu-to-llo-ods.md

Pass-arg flag工厂
lower-to-llomlir::tpu::createLowerToLLOPass (getArgument() @ 0x112041c0)
lower-to-mlomlir::tpu::createLowerToMloPass
tpu-tiling-propagation / tiling-propagationcreateTilingPropagationPass (getArgument() @ 0x132e0d80)
tpu-apply-vector-layoutcreateApplyVectorLayoutPass
tpu-infer-vector-layout / infer-vector-layoutcreateInferVectorLayoutPass
tpu-infer-memref-layoutcreateInferMemRefLayoutPass
tpu-relayout-insertion / relayout-insertioncreateRelayoutInsertionPass
tpu-canonicalize-{memory-space,mosaic,operations}createCanonicalize*Pass
(无 flag)createConvertIntegerMemrefsPass, createDebugAssertInsertionPass, createLogicalToPhysicalDeviceIdPass, createPreCanonicalizationOptimizationPasstpu 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 checksAddStreamBoundChecks, IndirectStreamBoundsCheck, LoadStoreBoundsCheck, InsertRunIdCheck, LlvmIntToPtrSafety, SyncFlagVerifier
Tile outlining / structureTileTaskOutlining, SinkOpsIntoTileTask, ExpandTiledMemRefs, TileTaskArgumentsSpill, InsertClearTileIbuf
AllocationAllocationAssignment, GlobalAllocationAssignment, TileOverlayAllocation, LegalizeHbmGlobals, PrepareHbmSpill
Scheduling / prefetchScheduleTileTaskPrefetch, LowerPrefetchTileTask, DelaysInsertion, TileBarrierInsertion
LoweringLowerToSparseCoreLlvm, LowerDynamicDimensionSize, LowerSequencerFunctions, PackedOperandsLowering, AlgebraicDecomposer
MiscHardwareTracing, SflagWaitInstrumentation, EmitTileOverlayHandler, ScsOverlaysTrampoline, MemrefAliasScopeAnnotation, PackUnpackCleaner, VectorUnroll, LoopUnroll, SparseCoreCanonicalizer, TileHostTransfer, LegalizeMegaChipModule
Pass-arg flagslegalize-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.mdllvmtpu-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
Schedulingtpu-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-passestpu-ssa-super-pass, tpu-enable-pipeliner-super-pass, tpu-enable-overlayer-passes
Bundle packing / encodingtpu-bundle-packer[-maintain-order/-maintain-topologial-order], tpu-encode-mcinst-bundles, tpu-post-bundle-lower, tpu-print-{mcinst,sc-mcinst}-encodings
Register alloc / spilltpu-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 / loweringtpu-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
Verifierstpu-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-bridgetpu-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-inlinertpu-instruction-fusiontpu-int2-auto-up-down-castertpu-tiling-propagationtpu-partition-assignment 既作为 HLO name() 字符串(上方)出现,也作为这里的 tpu-* flag 参数出现。它们是同一个逻辑 pass 暴露在两个注册层级上。把 tpu-* flag 和 HLO name() 字符串算作互不相交集合会重复计数;应将 HLO-level pass 视为规范名称,flag 视为它的 CLI alias。


Pipeline-Stage 名称 {#pipeline-stage-names}

两套阶段分类并存。Monolithic compile 用首字母大写字符串(传给 pipeline 构造函数的 std::string)标记其 HloPassPipeline 阶段;separate compiletpu_phase_functions.cc)用 phaseN_* 指标 token 标记其阶段。它们不能互换。两者都记录于 compile-phases.md

Monolithic HloPassPipeline Stage Strings

Stage stringFile-off角色
pre-optimization0x85ea920PreOptimizationPipeline(Phase 1)
post-optimization0x85ea85ePostOptimizationPipeline
Pre layout assignment0x84ef7bfpre-layout HLO
Layout assignment / XLA::JF Layout Assignment0x84ef7d5 / 0x84efbdfLayoutAssignment
Post layout assignment0x84ef6dbpost-layout HLO
HLO passes through layout assignment0x84ef757HloOptimizeThroughLayoutAssignment
HLO passes after layout assignment0x84ef6f2HloOptimizeAfterLayoutAssignment
Pre main fusion0x10963b16PostLayoutPreFusionHloOptimize
Main fusion / XLA::TPU main fusion pass— / 0x8522e3eTpuInstructionFusion bucket
Post main fusion / JF after main fusion0x85f7605 / 0x85f7616PostMainFusionHloOptimize
final_scheduler(SSO strcpy)基础 memory-schedule pipeline(Phase 7)
async_scheduling0x865dcabLHS + async-overlap pipeline(Phase 7)
final_scheduled0x8713d6apost-scheduling marker
simplification after layout assignment0x84ef715post-layout simplify

Separate-Compilation Phase Metric Tokens

Multi-phase compile 用 phaseN_* token 命名其阶段(separate-compilation timing 表中的 metric key)。phase1_hlo_optsphase3_linking 已按字节确认为从 xla::TpuCompiler::RegisterAllPhases @ 0xf849ec0 引用的字面量。“TLP” = TPU-Level-Program。

Token角色
phase0_stablehlo_to_hloStableHLO → HLO import
phase1_hlo_optsHLO optimization(上方 HLO passes)
phase2a_tlp_loweringHLO → MLIR-tpu(TPU-Level-Program)lowering
phase2b_deduped_loweringdedup'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_dimper-phase transfer/shard metric fields

NOTE — phaseN_* token 是指标 key,不是 pass 名称:它们出现在 timing/transfer 表中,绝不会作为读者可以运行的 pass 出现。VLOGxla_dump 过滤器会匹配上方的 stage 字符串和前面的 pass 名称;除非读取 separate-compilation timing 路径,否则不会匹配 phaseN_*。锚定 phase1_hlo_opts/phase3_linkingRegisterAllPhases 符号是注册入口,不是逐阶段 pass 主体。


未确认 / 不在本页

诚实列出范围限制,便于重新实现者知道需要重新推导什么:

  • 约 300 个 HLO name() 字符串按约定推断(INFER),未按字节确认。 name() 主体在运行时构造字符串(static const 或 string-builder);只有约 25 个带短横线名称作为 .rodata 字面量存在(CERTAIN 行)。类符号是确定的;精确返回字符串遵循 CamelCasekebab 约定,但未逐 pass 定位。
  • 每个 pass 的添加顺序不在这里。 本页是名称索引;顺序归 hlo-pass-registry.md(Phase-1..6 序列)、scheduler placement 和 fusion-bucket order 所有。
  • flag → pass gating 不在这里。 哪个 xla_tpu_* flag 启用哪个 AddPass<T>RunHloPasses0x1093a420)中的 .text 分支,不是本页恢复的字符串。
  • 242 个 tpu-* flag 的 llvm::Pass 类符号未恢复:只恢复了 cl::opt flag 字符串。
  • per-TpuVersion(jxc/pxc/vxc/glc/gfc)pass-set delta 未隔离:HAL 模块按版本命名,但 v3 与 v6e 哪些 pass 被跳过是 Target::tpu_version() 分支,本页未追踪。

交叉引用