HLO Pass Registry

地址、build-id 和符号名适用于 libtpu-0.0.40-cp314 wheel 中的 libtpu.so（build-id 89edbbe81c5b328a958fe628a9f2207d）。其他版本会有所不同；应将每个 VA 都视为绑定到特定版本。

摘要

TPU 后端上的每个 HLO 级变换都运行在一个容器类 xla::HloPassPipeline 和一个驱动方法 RunPassesInternal 中。不存在 TPU 私有的 pipeline 类型。TPU 编译器复用与 CPU 和 GPU 后端相同的 upstream-XLA HloPassPipeline，并由 Jellyfish pipeline-builder 函数实例化和填充（compile-phases.md 负责这些 builder 和 phase 顺序）。本页负责该顺序之下的机制：Run 驱动循环、每个 pass 实现的 HloPassInterface ABI、builder 调用的 AddPass/AddInvariantChecker 注册 API、由 flag 门控的 pass 禁用/启用过滤器（GetEnabledPasses），以及每个 mutating pass 后运行的 invariant-checker 重新验证循环。

熟悉 LLVM 的读者可以把它映射到 LLVM legacy PassManager，但差异是核心。HloPassPipeline 本身也是一个 HloPassInterface，所以 pipeline 可以像 pass 一样嵌套（HloPassFix<HloPassPipeline> 会把整个 sub-pipeline 重新运行到 fixed point）。Pass 不声明 analysis dependency，也没有 analysis cache；每个 pass 重新推导自己需要的信息，并返回一个 StatusOr<bool>，也就是“我是否改变了模块”这一位。Invariant checker 是第二个 pass 列表，运行相同的 Run ABI，但必须不修改，并且会在每个报告发生改变的 pass 之后重新验证，而不是只在 pipeline 结尾运行一次。整个驱动器写得很防御：它会用模块的独立 hash 交叉检查每个 pass 自己报告的 change bit；如果 pass 谎报是否修改过模块，编译器会 CHECK 失败。

本页是 ABI 与驱动器参考。哪些 pass 被添加、以什么顺序添加的枚举见 compile-phases.md（十阶段主干）和 hlo-pre-passes.md（前端 pre-pass 表）；372 项 name() RTTI 表面（_ZNK3xla*4nameEv 符号）在 hlo-pre-passes.md 中枚举；各个 pass 算法位于各自页面（algebraic-simplifier.md、layout-assignment.md、fusion-patterns.md、msa-overview.md）。这里仅记录容器、接口以及 add/gate/check 机制。

对重新实现而言，契约是：

HloPassInterface ABI。 每个 pass 和每个 checker 都是一个 vtable，其中有返回 string_view 的 name() slot（vtable+16）、返回 StatusOr<bool> 的 Run(module, execution_threads) slot，以及 is_pass_pipeline() 判别器。Pipeline 只是一个其 Run 为 RunPassesInternal 的 pass。
双列表容器。 HloPassPipeline 持有一个 passes_ vector 和一个独立的 invariant_checkers_ vector，外加一个 run_called_ latch。AddPass<T> 追加到前者；AddInvariantChecker<T> 追加到后者；如果在 Run 开始后调用，两者都会 CHECK 失败。
驱动循环。 RunPassesInternal 先运行一次 GetEnabledPasses，然后在 "pipeline-start" 运行所有 checker，随后对每个启用的 pass：计时、运行、用 HLO hash 验证其自报告 change bit，并在其发生改变时重新运行所有 checker。
flag gate。 GetEnabledPasses 会根据 --xla_disable_hlo_passes 和 --xla_enable_hlo_passes_only DebugOptions 列表过滤 passes_（两者互斥，由 CHECK 强制），过滤键是每个 pass 的 name()。


容器类	`xla::HloPassPipeline`（upstream XLA，非 TPU 私有）
源文件	`third_party/tensorflow/compiler/xla/hlo/pass/hlo_pass_pipeline.{cc,h}`
驱动器	`HloPassPipeline::RunPassesInternal<HloModule*>` @ `0x1c83ddc0`（反编译 786 行）
flag 过滤器	`HloPassPipeline::GetEnabledPasses` @ `0x1c83d0a0`
checker 循环	`HloPassPipeline::RunInvariantCheckers<HloModule*>` @ `0x1c840500`
HLO dump hook	`HloPassPipeline::MaybeDumpHloAndSaveFilenames` @ `0x1c83da00`
添加（checker）	`HloPassPipeline::AddInvariantChecker<HloCycleDetection>` @ `0x109459a0`; `<HloVerifier, TargetVerifierMetadata>` @ `0x1306d420`
Fixed-point wrapper	`HloPassFix<HloPassPipeline>::RunToFixPoint` @ `0x10955dc0`; `RunImpl` @ `0x10955d20`
Pass 接口	`xla::HloPassInterface` — `name()`（vtable+16）、`Run()`（`xla::HloPassInterface::Run` thunk）、`is_pass_pipeline()`
`passes_` vector	pipeline 对象内 offset +40（ptr）/ +48（count）
`invariant_checkers_` vector	offset +56（ptr）/ +64（count）
`run_called_` latch	offset +80（byte）
`name()` 符号	372 个 `_ZNK3xla4nameEv` 符号 — 137 个 `xla::jellyfish::`，24 个 `xla::tpu::sparse_core::`，211 个其他 `xla::`（并非全是 HLO pass；例如 `xla::Executable::name()` 也在集合中）
置信度	除非某行或标注另有说明，否则为 CONFIRMED（byte-anchored）

`HloPassInterface` ABI

目的

HloPassInterface 是每个 HLO pass 派生自的单一抽象基类。它是驱动器分发所依赖的契约：驱动器从不知道 pass 的具体类型，只知道其 vtable。TPU 后端在 211 个其他 xla::* 类之上添加了 137 个 xla::jellyfish::* 和 24 个 xla::tpu::sparse_core::* 的 name() override 类（总计 372 个 _ZNK3xla*4nameEv 符号；并非全是 HLO pass，该集合也包含 xla::Executable 等非 pass name() override），但每个具体 pass 都共享这个 vtable 形状，这正是 RunPassesInternal 循环可以类型无关的原因。

Vtable 布局

对驱动器重要的有三个 slot，它们从 RunPassesInternal、RunInvariantCheckers 和 GetEnabledPasses 的调用点恢复：

Slot	Offset	Signature	使用者
`name()`	vtable+16	`string_view name() const`	`GetEnabledPasses`（过滤键）、驱动器 VLOG、dump 文件名
`Run()`	通过 `xla::HloPassInterface::Run` 分发	`StatusOr<bool> Run(HloModule*, const flat_hash_set<string_view>& execution_threads)`	驱动器逐 pass 调用、checker 调用
`is_pass_pipeline()`	vtable+32	`bool is_pass_pipeline() const`	区分 pipeline-as-pass 和 leaf pass；vtable+32 slot 在 `RunPassesInternal` 的 change-verification 路径上未被执行（hash recording 由 `DebugOptions+928\|+911` 门控，不由此 slot 门控）[Confidence: MEDIUM on the +32 offset]

name() slot 是整个表面的逆向工程锚点：372 个 _ZNK3xla*4nameEv 符号几乎完整枚举了 override name() 的 XLA 类，因为每个具体 pass 都 override name()。不过该集合也捕获少量非 pass 的 name() override（xla::Executable、PjRt executable），因此它是 pass 列表的超集，而非精确计数。Pass 返回的字符串是它的 dashed-lowercase 名称（"tpu-int2-auto-up-down-caster"、"sharding-propagation"、"legalize-scheduling-annotations"，均 byte-confirmed），也就是 --xla_disable_hlo_passes flag 匹配的键。

NOTE — Run 的 boolean 是 pass 返回的唯一信号。 Run 返回 StatusOr<bool>：非 OK Status 终止 pipeline；OK true 表示“我改变了模块”，OK false 表示“未改变”。没有更丰富的结果：没有修改过的 computation 列表，也没有 invalidated-analysis 集合。驱动器的整个 change-tracking 都建立在这一位上，这也正是它会用模块 hash 独立审计该位的原因（见 §驱动循环）。

Pipeline 也是 pass

HloPassPipeline 本身派生自 HloPassInterface。它的 Run 是 RunPassesInternal；它的 is_pass_pipeline() 返回 true。这是 pipeline 可以嵌套的结构性原因：builder 可以 AddPass<HloPassPipeline> 一个已经填充好的 sub-pipeline，外层驱动器会把它当成一个 pass 运行。HloPassFix<HloPassPipeline> 实例（RunToFixPoint @ 0x10955dc0, RunOnChangedComputations @ 0x10955b00）证实了这一点：整个 sub-pipeline 会像单个 pass 一样被重新运行到 fixed point。

text

HloPassInterface (abstract)
  ├─ name() const                      → string_view   (vtable+16)
  ├─ Run(module, threads)              → StatusOr<bool> (the change bit)
  └─ is_pass_pipeline() const          → bool           (vtable+32)
        │
        ├─ TpuAlgebraicSimplifier, TpuInstructionFusion, …  (322 leaf passes)
        ├─ HloVerifier / HloCycleDetection / LegalizeSchedulingAnnotations  (checkers)
        └─ HloPassPipeline             ── is_pass_pipeline()==true, Run==RunPassesInternal
              └─ HloPassFix<HloPassPipeline>  ── re-runs the sub-pipeline to fixpoint

Pipeline 容器

目的

HloPassPipeline 是 builder 函数填充的对象。它是一个具名的双有序列表容器：pass 列表和 invariant checker 列表，外加 run_called_ latch，一旦执行开始就冻结两个列表。该容器不做调度，也不做 dependency analysis；顺序完全是插入顺序。

对象布局

从驱动器和 AddInvariantChecker 实例访问的字段 offset 恢复：

字段	Offset	类型	含义
vtable	+0	`HloPassInterface` vptr	pipeline-as-pass 分发
`name_`	+8…+32	内联 `std::string`（SSO）	pipeline 名称，例如 `"pre-optimization"`、`"Layout assignment"`、`"async_scheduling"`（均 byte-confirmed）
`passes_`	+40 / +48	`vector<unique_ptr<HloPassInterface>>` ptr / count	有序 pass 列表，由 `GetEnabledPasses` 读取（作为 `v15+4`/`v15+5`）
`invariant_checkers_`	+56 / +64	`vector<unique_ptr<HloPassInterface>>` ptr / count	checker 列表，由 `RunInvariantCheckers` 读取（作为 `a1+56`/`a1+64`）并由 `AddInvariantChecker` 追加
`run_called_`	+80	`bool`	latch — 在 `Run` 入口设为 true；之后的 `AddPass`/`AddInvariantChecker` 会 `CHECK` 失败

Pipeline 名称会出现在 FATAL 诊断中（"Pass '%s' in pipeline '%s' …"）。大小写混合的名称（"Layout assignment"、"pre-optimization"、"async_scheduling"）是 Jellyfish builder 传入 HloPassPipeline 构造函数的字符串；compile-phases.md 将每个名称映射到其 stage。

QUIRK — checker 列表与 pass 列表分离，并以不同 cadence 运行。 朴素的重新实现会把 verifier 作为普通 pass 合并进 pass 序列。libtpu 把 invariant_checkers_ 保存在第二个 vector 中，并在第一个 pass 前以及每个 changing pass 后运行整个 checker 列表，这形成对 checker 数量二次、对 pass 数量线性的验证调度。把它们合并进 passes_ 只会让每个 checker 在自己的 slot 运行一次，而不是在其他每个 pass 后运行，从而漏掉连续调度能捕获的格式错误中间态。

Add / Gate 机制

`AddPass` 和 `AddInvariantChecker`

两者都是 builder 函数在构造期间调用的模板。AddPass<T>(args...) 构造一个 T 并追加到 passes_（offset +40）；AddInvariantChecker<T>(args...) 构造一个 T 并追加到 invariant_checkers_（offset +56）。两者都由 run_called_ latch 进行 CHECK 防护。

AddInvariantChecker<HloCycleDetection> 实例（0x109459a0）端到端展示了该机制：

function AddInvariantChecker<T>(pipeline, ctor_args...):     // 0x109459a0 (HloCycleDetection)
    if pipeline->run_called_ == 1:                           // offset +80
        FATAL("!run_called_",                                // hlo_pass_pipeline.h:76
              "AddInvariantChecker cannot be called after Run")
    checker = new T(ctor_args...)                            // e.g. HloCycleDetection (vtable off_217F4558)
    pipeline->invariant_checkers_.push_back(                 // offset +56
        unique_ptr<HloPassInterface>(checker))
    return checker

AddPass<T> 针对 passes_（offset +40）具有相同形状；run_called_ CHECK 是共享的。该 latch 使 pipeline 一旦运行就不可变：不存在运行期间添加。

恢复出的 AddInvariantChecker 实例精确命名了 TPU pipeline 注册的三种 checker 类型（与 compile-phases.md 中的 MaybeAddInvariantCheckers 交叉引用）：

Checker 类型	Ctor metadata	VA	角色
`HloVerifier`	`unique_ptr<TargetVerifierMetadata>`	`0x1306d420`	带 TPU verifier metadata 的结构性 HLO verifier
`HloVerifier`	`unique_ptr<CpuGpuVerifierMetadata>` + name `"…"`	`0x14bcb340`	开源 CPU/GPU-metadata 形式（存在但不是 TPU 路径）
`HloCycleDetection`	(none)	`0x109459a0`	拒绝 HLO 图中的控制/数据 cycle
`LegalizeSchedulingAnnotations`	`Config&`	`0x10945740`	scheduling-annotation 合法性，注册为 checker

QUIRK — 同一个类既是 pass 又是 checker。 LegalizeSchedulingAnnotations（name() = "legalize-scheduling-annotations"）既出现在 pass 列表中，也作为 invariant checker 出现。作为 pass，它合法化 annotation（会修改）；作为 checker，它必须只读运行。接口不区分两者，都会调用 Run，所以“checker 不得修改”的规则是一种约定，驱动器只通过验证整个 checker sweep 前后模块 hash 不变来间接强制，而不是逐 checker 强制。

`AddInvariantChecker<HloVerifier>` — TPU verifier 接线

TPU-metadata 实例（0x1306d420）构造一个 HloVerifier（vtable off_21D2A9A0），其中持有从 unique_ptr 参数 move 进来的 TargetVerifierMetadata，并将 verifier 的 name 字段初始化为 "Unknown"。TargetVerifierMetadata 是 XLA VerifierMetadata 的 TPU 子类；开源默认值是 CpuGpuVerifierMetadata（0x14bcb340 实例）。重新实现必须提供 target-specific metadata 对象，让 verifier 强制 TPU shape/layout invariant，而不是通用 CPU/GPU 集合。

Flag gate — `GetEnabledPasses`

GetEnabledPasses (0x1c83d0a0) 在 RunPassesInternal 顶部运行一次。它基于每个 pass 的 name()，把完整 passes_ 列表转为驱动器实际运行的过滤后列表：

function GetEnabledPasses(pipeline, debug_options) -> vector<HloPassInterface*>:  // 0x1c83d0a0
    if debug_options.fast_path_flag(+553) == 1:            // short-circuit: all passes enabled
        return all_of(pipeline->passes_)
    disabled = set(debug_options.xla_disable_hlo_passes)   // repeated string @ DebugOptions+64
    enabled  = set(debug_options.xla_enable_hlo_passes_only)// repeated string @ DebugOptions+112
    FATAL_IF(!disabled.empty() && !enabled.empty(),        // hlo_pass_pipeline.cc:247
             "disabled_pass_names.empty() || enabled_pass_names.empty()")

    // (1) whole-pipeline gate: the pipeline's own name() can be disabled/enabled
    if disabled.contains(pipeline->name()):                // VLOG "Disable the full pass:"  cc:251
        return {}                                          // empty → pipeline runs nothing
    if !enabled.empty() && !enabled.contains(pipeline->name()):  // cc:256
        return {}

    // (2) per-pass gate: keep a pass iff its name() is not in the disable set
    out = []
    for pass in pipeline->passes_:                         // ptr/count @ pipeline+40 / +48
        if !disabled.contains(pass->name()):
            out.push_back(pass)
    return out

GOTCHA — disable 和 enable-only 互斥，且禁用一个 pipeline 名称会静默整个 stage。 有两个陷阱。第一，同时设置 --xla_disable_hlo_passes 和 --xla_enable_hlo_passes_only 是硬 CHECK 失败（cc:247），不是合并；允许两者共存的重新实现会立刻偏离。第二，disable/enable 检查首先针对 pipeline 自身的名称运行：因为每个嵌套 pipeline（"Layout assignment"、"async_scheduling"）本身也是具名 pass，禁用该名称会返回空 pass 列表并跳过整个 stage，而不只是跳过某个 pass。Byte-confirmed 的 VLOG 字符串 "Passes disabled by --xla_disable_hlo_passes: "（cc:238）、"Passes enabled by --xla_enable_hlo_passes_only: "（cc:243）、"Disable the full pass: "（cc:251）和 "Enable the full pass: "（cc:256）锚定了每个分支。
NOTE — 这是基于名称的 gating，不是 xla_jf_*/xla_tpu_* feature flag。 GetEnabledPasses 只尊重两个 debug-options pass-name 列表。启用或关闭 pass 的逐 feature flag（xla_jf_*、xla_tpu_*、xla_msa_*）在每个 builder 函数内部读取：if (env->flag) AddPass<T>(...)，发生在 pass 进入 pipeline 之前。因此，被 flag 禁用的 pass 根本不会添加；被名称禁用的 pass 会先添加，然后在这里过滤掉。这两种机制可以组合；compile-phases.md 负责 builder-level flag gating，本页负责 name-level filter。flag→pass 绑定谓词没有逐一追踪。[Confidence: HIGH on the name filter; the feature-flag gating is documented on the builder pages.]

驱动循环

目的

RunPassesInternal (0x1c83ddc0) 是 registry 的核心：它按顺序单次遍历启用的 pass，运行每个 pass 并重新验证 invariant。这里也是编译器防御性 change-auditing 所在，也就是最可能让重新实现者意外的部分。

算法

function RunPassesInternal(pipeline, debug_options, module, threads) -> StatusOr<bool>:  // 0x1c83ddc0
    enabled = GetEnabledPasses(debug_options)              // the filtered pass list
    PushAnnotation("pipeline: " + pipeline->name())        // tsl::profiler scope

    // (A) initial invariant sweep
    status = RunInvariantCheckers("pipeline-start")        // cc:151 on failure
    if !status.ok(): return status
    RecordPassStartMetadata("pipeline-start")
    MaybeDumpHloAndSaveFilenames("pipeline-start", next_pass_name)
    RecordPassEndMetadata()

    // verification mode: read DebugOptions hash-check toggles once
    check_hash = module.flags(+928)                        // "verify the pass changed the HLO"
|| module.flags(+911)                        //   (combined into v139)
    changed_overall = false

    for i, pass in enumerate(enabled):                     // passes+40 / count+48
        timer = ScopedLoggingTimer("HLO pass: " + pass->name())  // cc lambda#2
        VLOG(1) "  HLO pass " << pass->name()              // cc:174
        VLOG(2) "  Number of instructions: " << module.instruction_count()  // cc:180

        if check_hash:
            hash_before = AbslHashValue(module)            // independent fingerprint

        PushAnnotation("pass: " + pass->name())
        RecordPassStartMetadata(pass->name())
        status = pass->Run(module, threads)                // the only mutation point
        if status.ok():                                    // status carried the change bit
            module.Cleanup()                               // drop dead state post-pass
        else:
            return status.AddSourceLocation(cc:191)        // pass failed → abort
        changed = status.value()                           // the self-reported bit

        // (B) audit the change bit against the hash  (cc:116 / cc:125)
        if check_hash:
            hash_after = AbslHashValue(module)
            if hash_after != hash_before && !changed && module.flags(+928):
                FATAL("Pass '%s' in pipeline '%s' reported that it did not "  // cc:116
                      "change the HLO but the hash of HLO was changed …", …)
            if hash_after == hash_before && changed && module.flags(+911):
                FATAL("Pass '%s' in pipeline '%s' reported that it changed "  // cc:125
                      "the HLO but the hash of HLO was not updated …", …)

        MaybeDumpHloAndSaveFilenames(pass->name(), next_pass_name)  // dump after this pass
        RecordPassEndMetadata(changed)
        changed_overall |= changed

        // (C) re-validate invariants only if this pass actually changed the module
        if changed:
            VLOG(3) "  Pass caused changes " << pass->name()       // cc:206
            status = RunInvariantCheckers(pass->name())            // cc:213 on failure
            if !status.ok(): return status.AddSourceLocation(cc:213)

    return changed_overall

驱动器保证什么

一次一个 pass，插入顺序，不重排。 enabled 从前到后遍历；没有 priority，没有 dependency graph，没有 analysis cache。为“效率”而重排的重新实现会生成不同程序。
每个成功 pass 后执行 module.Cleanup()。 Pass 留下的 dead computation/instruction 会在下一个 pass 看到模块前被丢弃；因此 pass 不必完美清理，但下一个 pass 总会看到清理后的模块。
Checker 只在changing pass 后重新运行。 返回 false（无变化）的 pass 不会触发 checker sweep，模块被假定仍然有效。这就是 compile-phases.md 标注所指的 cadence：每个mutating pass 后验证，而不是在 phase 边界验证。
Change bit 会被审计，而不是被信任。 两条 FATAL 路径（cc:116、cc:125）捕获 pass 撒谎的两个方向：改了但说没改，以及说改了但没改。两者分别由 DebugOptions toggle（module+928、module+911）门控，所以审计是 opt-in（debug/CI 构建），但代码始终存在。

GOTCHA — 启用时，change-bit 审计会对整个模块每个 pass hash 两次。 hash-check toggle 打开后，每个 pass 都要付出两次完整 AbslHashValue(HloModule) 遍历，以及任何改变后的一次 RunInvariantCheckers sweep。这是 debug/verification 模式，不是 production 路径；把审计接入 hot path 的重新实现会看到大型模块编译时间膨胀。审计存在是为了捕获最糟糕的 pass bug（陈旧 change bit 静默禁用下游 fixed-point convergence），所以正确做法是保留它，并保持门控。

`RunInvariantCheckers` — checker sweep

function RunInvariantCheckers(module, after_pass_name, threads) -> Status:  // 0x1c840500
    if invariant_checkers_.empty(): return OK              // a1+64 == 0 fast path
    for checker in invariant_checkers_:                    // a1+56 / a1+64
        VLOG(1) "    Invariant checker " << checker->name()  // cc:83
        status = checker->Run(module, threads)             // SAME ABI as a pass
        if status.ok():
            module.Cleanup()                               // checkers may schedule cleanup
        else:
            return status.AddSourceLocation(hlo_pass_pipeline.h:147)
    return OK

Checker 通过与 pass 相同的 HloPassInterface::Run slot 分发；唯一差异是策略（它们不得修改），以及任何 checker 返回非 OK status 都会中止整个 pipeline。每次调用 sweep 时都会运行完整 checker 列表。

Fixed-Point 包装 — `HloPassFix`

某些 pass 和整个 sub-pipeline 必须运行到收敛，而不是只运行一次。HloPassFix<P> 包装一个 pass P（或一个 pipeline），并重复运行它，直到它报告无变化。恢复出的实例证实该机制对 leaf pass 和 pipeline 都是泛型的：

Wrapper	`RunToFixPoint` VA	重新运行的内容
`HloPassFix<HloPassPipeline>`	`0x10955dc0`	整个 sub-pipeline 到 fixed point
`HloPassFix<HloDCE>`	`0x1d6d7a60`	dead-code elimination 到 fixed point
`HloPassFix<jellyfish::TpuReduceWindowRewriter>`	`0x109589e0`	TPU reduce-window rewrite
`HloPassFix<jellyfish::MosaicFusion>`	`0x1095e800`	Mosaic-kernel fusion
`HloPassFix<AllReduceReassociate>`	`0x1095f5e0`	collective reassociation
`HloPassFix<ReduceWindowRewriter>`	`0x14bd0980`	OSS reduce-window
`HloPassFix<ReduceScatterReassociate>`	`0x109603c0`	collective reduce-scatter reassoc
`HloPassFix<WhileLoopConstantSinking>`	`0x12ef1280`	while-loop constant sinking
`HloPassFix<AllReduceReduceScatterReorder>`	`0x109611a0`	collective reorder

HloPassFix 暴露相同的 HloPassInterface ABI（RunImpl / RunOnChangedComputations / RunOnChangedComputationsOnce），因此外层驱动器把 fixed-point wrapper 当成普通 pass 处理，而它返回的 change bit 是所有迭代的 OR。

NOTE — 非收敛行为。 Upstream XLA 会限制 HloPassFix 迭代次数，并在某个 flag 下，如果 pass 震荡而非收敛则 fatal。采样字符串中未定位到 TPU crash-on-non-convergence 行为的 flag 字符串；HloPassFix fixed-point 机制本身是 CONFIRMED（上面的 RunToFixPoint/RunOnChangedComputations vtable 三元组），但 crash-on-divergence 声明为 [Confidence: LOW]。见 compile-phases.md 上的对应标注。

Builder 如何使用 Registry

Jellyfish pipeline-builder 函数（compile-phases.md）是 AddPass/AddInvariantChecker 的唯一调用者。结构上恢复出的构造模式是：

text

CreateHloPipeline (0x1093efe0)                 ── allocates the top HloPassPipeline
  ├─ AddInvariantChecker<HloVerifier(TargetVerifierMetadata)>   (via MaybeAddInvariantCheckers 0x10944600)
  ├─ AddInvariantChecker<HloCycleDetection>
  ├─ AddInvariantChecker<LegalizeSchedulingAnnotations>
  ├─ AddPass<...> × N             ── the pre-opt / sharding / layout / fusion passes, flag-gated
  └─ AddPass<HloPassPipeline>     ── nested named sub-pipelines ("Layout assignment", "async_scheduling", …)
        └─ (each nested pipeline re-runs MaybeAddInvariantCheckers at its own head)

MaybeAddInvariantCheckers (0x10944600) 是在每个嵌套 pipeline 头部注册三个 checker 的唯一 helper，这就是连续重新验证在每个嵌套层级都成立、而不只是顶层成立的原因。Builder 返回后，填充好的 pipeline 会被 RunPassesInternal 精确调用一次，随后 run_called_ latch 会拒绝任何迟到的修改。

QUIRK — registry 是每次编译新建的，不是静态表。 尽管使用了“registry”表述，但这里没有像 LLVM PassRegistry 那样的全局 pass registry。每次 compile 都会构造新的 HloPassPipeline，builder 命令式地向其中 AddPass（带 flag gating），它运行一次，然后销毁。“registry”是 372 个 name() 的类型表面（hlo-pre-passes.md）加上这个逐编译实例化模式，而不是持久注册 DAG。（确实持久存在的 google_init_module_* DAG 注册的是 emitter 和 HAL factory，不是 HLO pass。）

置信度摘要

声明	证据
`HloPassPipeline::RunPassesInternal` 是 HLO 驱动器	反编译 `0x1c83ddc0`，源文件 `hlo_pass_pipeline.cc`
Pass/checker 共享 `HloPassInterface::Run`（`StatusOr<bool>`）	`RunPassesInternal` 和 `RunInvariantCheckers` 都调用 `xla::HloPassInterface::Run`
双列表：`passes_`(+40/+48)、`invariant_checkers_`(+56/+64)、`run_called_`(+80)	driver/`GetEnabledPasses`/`AddInvariantChecker` 读取的 offset
`AddInvariantChecker` 在 `Run` 后 `CHECK` 失败	`0x109459a0`、`0x1306d420` 中 `hlo_pass_pipeline.h:76` 的 FATAL `"!run_called_"`
三个 TPU checker：`HloVerifier(TargetVerifierMetadata)`、`HloCycleDetection`、`LegalizeSchedulingAnnotations`	listed VA 处的 `AddInvariantChecker<…>` 实例
`GetEnabledPasses` 根据 `xla_disable_hlo_passes`/`xla_enable_hlo_passes_only` 过滤，且两者互斥	反编译 `0x1c83d0a0`，byte 字符串 `cc:238/243/247/251/256`
驱动器审计 change bit 与 HLO hash，不匹配则 FATAL	`cc:116`/`cc:125` FATAL 字符串 + `AbslHashValue(HloModule)` 调用
Checker 在每个 changing pass 后重新运行（不只一次）	循环内 `if changed: RunInvariantCheckers` @ `cc:213`
`HloPassFix<HloPassPipeline>` 把 sub-pipeline 重新运行到 fixed point	`RunToFixPoint` `0x10955dc0` + leaf-pass `HloPassFix` 实例
`HloPassFix` 的 crash-on-non-convergence flag	采样字符串中未定位到 flag 字符串
Builder 是唯一的 `AddPass`/`AddInvariantChecker` 调用者；逐编译实例化	`CreateHloPipeline`/`MaybeAddInvariantCheckers` 调用结构

交叉引用

Compile Phases 0–3 — 十阶段主干，以及填充此容器的 Jellyfish builder 函数（CreateHloPipeline、PreOptimizationPipeline、MaybeAddInvariantCheckers）。
The TPU Compiler — Part V 导向；HLO pass pipeline（Family 1）在五个编译阶段和 IR 层栈中的位置。
HLO Pre-Passes — 此驱动器运行的枚举 pass 目录（372 项 name() RTTI 表面），以及逐 pass HLO invariant。
Algebraic Simplifier — 实现 HloPassInterface ABI 的代表性 leaf pass（TpuAlgebraicSimplifier）。
Layout Assignment — 作为 HloPassInterface 在 through-layout pipeline 内运行的 codegen-gating analysis。
Fusion Patterns — TpuInstructionFusion，Phase 5 fusion stage 中添加的主 fusion pass。
MSA Overview — memory-space assignment，由 RunMemorySpaceAssignment 在 post-pipeline 驱动。
返回索引
二进制： extracted/libtpu-0.0.40-cp314-cp314-manylinux_2_31_x86_64/libtpu/libtpu.so（build-id 89edbbe81c5b328a958fe628a9f2207d）

HLO Pass Registry ​

摘要 ​

HloPassInterface ABI ​

目的 ​

Vtable 布局 ​

Pipeline 也是 pass ​

Pipeline 容器 ​

目的 ​

对象布局 ​

Add / Gate 机制 ​

AddPass 和 AddInvariantChecker ​

AddInvariantChecker<HloVerifier> — TPU verifier 接线 ​

Flag gate — GetEnabledPasses ​

驱动循环 ​

目的 ​

算法 ​

驱动器保证什么 ​

RunInvariantCheckers — checker sweep ​

Fixed-Point 包装 — HloPassFix ​

Builder 如何使用 Registry ​

置信度摘要 ​

交叉引用 ​

HLO Pass Registry

摘要

`HloPassInterface` ABI

目的

Vtable 布局

Pipeline 也是 pass

Pipeline 容器

目的

对象布局

Add / Gate 机制

`AddPass` 和 `AddInvariantChecker`

`AddInvariantChecker<HloVerifier>` — TPU verifier 接线

Flag gate — `GetEnabledPasses`

驱动循环

目的

算法

驱动器保证什么

`RunInvariantCheckers` — checker sweep

Fixed-Point 包装 — `HloPassFix`

Builder 如何使用 Registry

置信度摘要

交叉引用