Skip to content

代数简化器

以下所有符号名称、VA 和构建 ID 适用于 libtpu-0.0.40-cp314 轮中的 libtpu.so(构建 ID 89edbbe81c5b328a958fe628a9f2207d、构建 libtpu_lts_20260413_b_RC00)。 .text VMA == 文件偏移量。其他版本有所不同;将每个 VA 视为版本固定。

摘要

代数简化器是 TPU 编译器的窥孔重写器 - HLO 遍历遍历每个计算并用更便宜的等效项替换本地可识别的指令模式:恒等消除(add x 0multiply x 1x AND true)、倒数折叠(当 b 为常数时为 a/b → a * (1/b))、min/max → clamp 强度降低、广播/重塑/转置规范化和折叠、点/卷积操作数转置吸收、通过 GetTupleElement 进行死元组修剪,以及算术规范化的长尾。在 TPU 后端,它是 HLO 预传递管道的 pass #60(请参阅 hlo-pre-passes.md),实例化不是开源 xla::AlgebraicSimplifier,而是 xla::jellyfish::TpuAlgebraicSimplifier — 一个子类,其访问者 (TpuAlgebraicSimplifierVisitor) 发出 17 个 Handle* 方法(16 个重写基本处理程序,其中包含TPU 布局感知变体,加上基础 AlgebraicSimplifierVisitor 提供的 56 个处理程序 之上的一个仅 TPU HandleRngBitGenerator)。

了解 LLVM 的读者应该将其视为 HLO 图的 InstCombine,有两个结构差异。首先,调度不是按利益排序的模式匹配(重写调度页面 中记录的 MLIR PatternApplicator 机制是不同 IR 上的单独引擎)。它是普通的 C++ 基于操作码的虚拟调度:基类 xla::DfsHloVisitorBase<HloInstruction*> 是一个约 140 个插槽的 vtable,每个 HloOpcode 一个 Handle<Op> 插槽,重写过程仅覆盖它关心的操作码的插槽,并且 HloInstruction::Visit 每条指令只发出一个间接调用。规则的“注册表”C++ 类层次结构。其次,简化器被包裹在运行到固定点循环中:重新访问每个计算,直到完整的传递没有改变为止,通过循环诊断最多进行 50 次迭代。

此页面记录了三件事,每件事都以字节为锚定:(1) 访问者调度 — 操作码 vtable 机制、两个覆盖集(基本 XLA 与 TPU 子类)以及每次计算 Run 条目; (2) 代表性重写规则目录 — 进行大量简化的处理程序,以及每个调用的辅助函数; (3) 运行到固定点驱动程序 TpuAlgebraicSimplifier::RunImplAlgebraicSimplifierOptions 门(启用/禁用规则系列的 TPU 特定旋钮)。它不会重新派生每个处理程序主体(所有 56 个基本处理程序的规则算术都是上游 XLA)、MLIR 端 PatternApplicator 或融合优先级队列 - 这些都有自己的页面。

TPU 通过等级xla::jellyfish::TpuAlgebraicSimplifier(xla::AlgebraicSimplifier 的子类)
TPU 运行至定点驱动程序TpuAlgebraicSimplifier::RunImpl(HloModule*, execution_threads) @ 0x13443f40
基本运行驱动程序xla::AlgebraicSimplifier::RunImpl(HloModule*, …) @ 0x1dd488c0
每个计算条目AlgebraicSimplifierVisitor::Run(HloComputation*, AlgebraicSimplifierOptions const&, AlgebraicSimplifier*) @ 0x1dd010e0bool (已更改)
TPU 访客xla::jellyfish::TpuAlgebraicSimplifierVisitor(虚表0x2190bcc8); 17 Handle*(16 个覆盖 + 1 个仅 TPU)
基本访客xla::AlgebraicSimplifierVisitor(虚表0x21d1d1e8,147个插槽); 56 Handle* 覆盖
调度基地xla::DfsHloVisitorBase<HloInstruction*>(虚函数表 0x21d2c320,140 个虚拟插槽 = 每个操作码一个 Handle<Op>)
选项结构xla::AlgebraicSimplifierOptions(dtor 0x1343ba00);内联存储在 +8 的通道中
通行证添加 (TPU)HloPassPipeline::AddPass<TpuAlgebraicSimplifier, Target const&, AlgebraicSimplifierOptions&> @ 0x10954400
通过添加(基础)HloPassPipeline::AddPass<AlgebraicSimplifier, AlgebraicSimplifierOptions&> @ 0x14bcb600
TPU源单元platforms/xla/service/jellyfish/tpu_algebraic_simplifier.cc(绳锚式,第 961/985/991 行)
基本源单元…/xla/hlo/transforms/simplifiers/algebraic_simplifier.cc(绳锚式)
固定点帽每次计算 50 次迭代 (0x32),然后是一个非致命的“循环简化循环”LOG
匹配EDSLxla::match:: 组合器(二进制范围内 1,117 个 Match<> 实例化)
信心已确认(字节锚定),除非行或标注另有说明

它在哪里

简化器在编译管道的第 1 阶段内运行,作为有序 HLO 预传递集 (hlo-pre-passes.md) 的传递 #60,位于 HloOptimizeThroughLayoutAssignment 的尾部(就在 HloPassFix<TpuReduceWindowRewriter> 定点之后、GatherOptimizer 之前)。它运行布局后分配,这就是 TPU 子类存在的全部原因:TPU 覆盖是布局感知 - 它们参考 IsValidLayout / UpdateLayout,因此重写永远不会产生 tpu→LLO 合法化器无法实现的平铺布局的张量。相比之下,开源简化器将布局视为不透明。

text
… LayoutAssignment (commits tiled layouts)
  → HloPassFix<TpuReduceWindowRewriter>
  → TpuAlgebraicSimplifier (#60)   ← THIS PAGE   RunImpl @ 0x13443f40
      └─ for each HloComputation (skip if >1 caller):
           AlgebraicSimplifierVisitor::Run(comp, options, this)   [vtable-dispatched]
           re-run to fixpoint (cap 50) if options.run_to_fixed_point()
  → GatherOptimizer → AllReduceSimplifier → …
```text

相同的 `RunImpl` 是传递执行的主体,无论它添加到何处; [hlo-pass-registry.md](hlo-pass-registry.md) 将通行证编入其 `HloPassInterface` RTTI 条目下,而 [optimization-barrier.md](optimization-barrier.md) 记录该访问者的一个特定处理程序(通过 `OptimizationBarrier` 进行死元组修剪),作为围绕 opt-barrier 进行操作的*唯一*重写。

---

## 调度:操作码键控虚拟表,而不是模式池

重新实现者的症结是:没有候选列表,没有效益分数,也没有每操作模式桶。规则 **选择** 是 C++ vtable。

### 基础虚函数表

`xla::DfsHloVisitorBase<HloInstruction*>`(虚表 `0x21d2c320`)是 **140 个虚拟插槽宽**(`0x21d2c790 − 0x21d2c320 = 0x470` 原始条目,减去顶部偏移量 + RTTI)。每个插槽都是一种 `Handle<Opcode>` 方法。槽遍历(解析 `R_X86_64_RELATIVE` 重定位加数)将它们标记为每个操作码一个:

```text
slot 0/1  ~DfsHloVisitorBase (D2/D0)
slot 2    HandleElementwiseUnary      slot 3    HandleElementwiseBinary
slot 4/5  __cxa_pure_virtual          (HandleClamp / HandleSelect — pure in base)
slot 6    HandleMaximum               slot 7    HandleMinimum
slot 9    HandleConvert               slot 10   HandleBitcastConvert
slot 11   HandleStochasticConvert     slot 12   HandleCopy
slot 13   HandleComplex               slot 14   HandleMultiply
slot 18   HandlePower  slot 19 HandleSqrt  slot 20 HandleRsqrt  slot 21 HandleCbrt
…                                       (140 total; __cxa_pure_virtual where the base
                                         forces the subclass to implement — e.g.
                                         HandleParameter / HandleConstant / HandleReduce)

具体的 HloInstruction 的操作码字节(此版本:*(BYTE*)(inst + 12),与 optimization-barrier.md 上记录的相同 +12 偏移量)对插槽进行索引。 HloInstruction::Visit(visitor) 读取操作码并通过该槽发出间接调用。 每条指令一个间接调用。 传递该子类 DfsHloVisitorWithDefault 为它们不覆盖的每个操作码获取一个无操作 DefaultAction,因此未处理的操作码只是保持不变。

为什么这不是 MLIR PatternApplicator

libtpu 在两个不同的 IR 上包含两个结构不同的重写调度引擎,重新实现者不得将它们混为一谈。 MLIR 端(stablehlo/tpu/llo 方言)使用 mlir::PatternApplicator — 一个 RewritePattern 对象池,每个对象携带一个 16 位 PatternBenefit,将每个 OperationName 存储到一个DenseMap,按效益降序排列稳定,派出首场比赛胜利。 HLO 代数简化器不使用

HLO 代数简化器(本页)MLIR PatternApplicator
IRxla::HloModule / HloInstructionMLIR Operation(tpu/llo/stablehlo 方言)
规则注册表C++ 类层次结构(覆盖 vtable 槽)FrozenRewritePatternSet RewritePattern*
选择操作码上的虚拟调度(每个操作码一个插槽)按操作种类排序的利益桶,首场比赛获胜
订购none — 每个操作码只有一个处理程序16 位 PatternBenefit 位于 Pattern+0x8,稳定排序的 DESC
匹配逻辑xla::match:: EDSL 在线评估 每个候选人matchAndRewrite,带有 canApply
运行循环重新访问固定点计算(上限 50)贪婪工作列表 深度感知转换合法化器

实际结果:如果不编辑访问者类,就无法为代数简化器“注册额外的规则”。 pass 执行的简化集在编译时由它覆盖的 Handle<Op> 插槽固定 - 这正是 TPU 专门化是子类(TpuAlgebraicSimplifierVisitor)而不是添加模式的原因。

两个倍率设置

重写过程的表面正是它覆盖的一组虚函数表槽。这里有两个覆盖集很重要,两者都直接从反编译中枚举(每个 Handle<Op> 都是其自己的发出函数)。

基础 — xla::AlgebraicSimplifierVisitor(虚函数表 0x21d1d1e8,147 个插槽):56 个覆盖。

text
Abs Add AllGather AllReduce AllReduceOrReduceScatter AllToAll And Bitcast
BitcastConvert Broadcast Clamp Compare Complex Concatenate Conditional Constant
Convert Convolution Copy CustomCall Divide Dot DynamicSlice DynamicUpdateSlice
Exp Gather GetTupleElement Imag Iota Log Map Maximum Minimum Multiply Negate Not
OptimizationBarrier Or Pad Power Real Reduce ReducePrecision ReduceScatter
ReduceWindow Remainder Reshape Reverse Rsqrt Scatter Select Slice Sort Sqrt
Subtract Transpose
```text

**TPU — `xla::jellyfish::TpuAlgebraicSimplifierVisitor`(vtable `0x2190bcc8`):17 `Handle*`** — 16 个覆盖基础处理程序 + 仅 TPU `HandleRngBitGenerator`(每个 TPU 处理程序采用 `Target const*`,以便可以查询每代布局模型):

```text
Add Bitcast Broadcast Compare Concatenate Convert Convolution Copy Dot
DynamicSlice DynamicUpdateSlice Pad Reshape RngBitGenerator Slice Sort Transpose

注意 - 对覆盖计数进行反编译交叉检查。 基本集是 56 不同的 AlgebraicSimplifierVisitor::Handle* 符号(顶级 Handle<Op>(HloInstruction*) 方法),TPU 集是 17 不同的 TpuAlgebraicSimplifierVisitor::Handle* 符号,通过列出发出的处理程序函数 (nm -C | rg) 进行枚举。在 17 个 TPU 处理程序中,有 16 个覆盖相同操作码的基本处理程序,其中一个 HandleRngBitGeneratorTPU-only:基本 AlgebraicSimplifierVisitor 不发出 HandleRngBitGenerator 符号,TPU 子类 添加 它 (0x13443440)。 RTTI vtable 宽度证实了 slot 计数(与覆盖计数无关):基础 vtable 0x21d1d1e8 跨越到其类型信息 0x21d1d690 = 0x4a8/8 = 149 条目(在偏移到顶部 + RTTI 后约为 147 个虚拟槽); 0x2190bcc8 处的 TPU 子类 vtable 具有相同的宽度。 [信心:已确认两个处理程序集和仅 TPU RngBitGenerator 事实。]

每个 TPU 重写都是一个改进:调度到达最派生的槽,因此对于 TPU 主体运行的 16 个 TPU 重写操作码(它可能委托给或替换基本逻辑),对于 RngBitGenerator 仅 TPU 处理程序运行,而对于其他 40 个基本处理程序,基本主体运行不变。每个处理程序“内部”的匹配逻辑是用 xla::match:: 组合器 EDSL(Op(),形状和操作数匹配器 - LLVM m_* 的 HLO 类似物)编写的,内联评估,而不是效益排名。


每次计算条目:AlgebraicSimplifierVisitor::Run

AlgebraicSimplifierVisitor::Run(comp, options, simplifier) (0x1dd010e0) 是由基础驱动程序和 TPU 驱动程序调用的单计算工作线程。它的签名,准确恢复:

cpp
// 0x1dd010e0
bool AlgebraicSimplifierVisitor::Run(
    HloComputation* computation,
    const AlgebraicSimplifierOptions& options,
    AlgebraicSimplifier* simplifier);   // returns: did anything change?
```text

它重置访问者的 `changed` 标志,按后序遍历计算指令,并通过操作码 vtable(上面)调度每条指令。每个执行重写的 `Handle<Op>` 都会回调简化器的替换助手(`ReplaceWith…`、`ReplaceInstruction`),这些助手设置 `changed` 标志并更新工作列表。 `Run` 返回该标志。 `true` 返回是驱动程序的固定点循环监视的内容。

---

## 运行至修复点驱动程序:`TpuAlgebraicSimplifier::RunImpl`

`TpuAlgebraicSimplifier::RunImpl`(`0x13443f40`)是通证体。恢复的结构(VLOG 转储之前/之后和迭代上限是字节精确的):

```cpp
// 0x13443f40  — name() anchored to tpu_algebraic_simplifier.cc
absl::StatusOr<bool> TpuAlgebraicSimplifier::RunImpl(
    HloModule* m,
    const absl::flat_hash_set<absl::string_view>& execution_threads) {
  // VLOG(2): "TpuAlgebraicSimplifier::RunImpl(), before:\n" + m->ToString()
  //          tpu_algebraic_simplifier.cc:961
  bool changed = false;
  for (HloComputation* comp : m->computations(execution_threads)) {
    // GATE: only simplify computations with <= 1 caller_instructions().
    //   v19 = comp->caller_instructions(...).size();  if (v19 > 1) skip.
    if (comp->caller_instructions(/*317*/).size() > 1) continue;

    if (AlgebraicSimplifierVisitor::Run(&visitor, comp, options_, this)) {
      changed = true;
      // RUN TO FIXPOINT — only if options_.run_to_fixed_point() (the bool at
      //   this + 135) is set; otherwise one pass per computation.
      if (run_to_fixed_point_ /* *(BYTE*)(this+135) == 1 */) {
        long runs = 1;
        while (runs - 1 < 0x32 /* 50 */) {            // hard cap
          ++runs;
          if (!AlgebraicSimplifierVisitor::Run(&visitor, comp, options_, this))
            break;                                    // converged
        }
        if (runs - 1 >= 0x32) {
          // tpu_algebraic_simplifier.cc:985 — NON-FATAL LOG, not a CHECK
          LOG(WARNING) << "Algebraic simplifier is likely stuck in a circular "
                          "simplification loop and ran for " << runs
                       << " runs on computation " << comp->name();
        }
      }
    }
  }
  // VLOG(2): after-dump at tpu_algebraic_simplifier.cc:991
  return changed;
}

重新实现者必须保留的三个事实,全部经过字节确认:

  • 单调用者门。 仅当 ≤ 1 caller_instructions() 时,计算才会被简化。此处跳过从多个站点调用的计算(它们通过其他通道以内联/扁平形式简化);在简化器的局部假设下,就地重写乘法调用的计算是不合理的。
  • Fixpoint 是选项门控的,上限为 50。 仅当 run_to_fixed_point 布尔值(this + 135 处的字节)为 1 时,才会触发重新运行循环。每次重新运行都是经过相同计算的完整 Run;收敛是“Run 返回false”。如果循环达到 50 次迭代而不收敛,它会发出 非致命 LOG (tpu_algebraic_simplifier.cc:985) 命名计算和运行计数,然后继续 - 它不会 CHECK 失败或中止。 (对比 hlo-pre-passes.md 上的管道级 HloPassFix 收敛行为,这是一个单独的机制。)
  • 选项位于通证内。 RunImplthis + 8 处读取 options_,在 this + 135 处读取 run_to_fixed_point 标志; AddPass 蹦床 (0x10954400) 将 AlgebraicSimplifierOptions 参数内联复制到从 +8 开始的 184 字节 (0xB8) 传递对象中。没有共享/全局选项单例——每个通道实例都有自己的门配置。

底座xla::AlgebraicSimplifier::RunImpl(0x1dd488c0)形状相同,不含TPU Target; TPU 子类的不同之处仅在于将调度路由到 TPU 访问者以及布局感知覆盖。


代表重写规则

规则机构主要是XLA的上游;本节对承担大部分工作的处理程序和每个调用的辅助函数进行了分类,因此重新实现者知道规则表面以及 TPU 专业化的分歧点。下面的每个地址都是此构建中的一个发出函数。

Base-XLA 处理程序(40 个在 TPU 上运行不变)

处理程序地址代表性简化(恢复表面)
HandleAdd(底座)x + 0 → x; x + (-y) → x - y;恒定折叠;常数的重新关联
HandleMultiply0x1dd...x * 1 → xx * 0 → broadcast(0)x * 2^k强度图案
HandleDivide0x1dd106a0a / b → a * (1/b)b不变时; x / x → 10 / x → 0;按 dtype 开关(IntegralTypeSwitch 0x1dd61e00FloatingPointTypeSwitch 0x1dd63d20)
HandleMaximum/HandleMinimum0x1dd22a40 / 0x1dd23960进给 MinMaxToClamp (0x1dd23380): max(min(x,hi),lo) → clamp(lo,x,hi)
HandleClamp / HandleSelect0x1dd23f00 / 0x1dd421e0冗余钳位消除; select(true,a,b) → a
HandleBroadcast0x1dd27a80广播崩溃;标量广播规范化;接收器按元素广播过去
HandleReshape0x1dd2eb40重塑重塑崩溃; reshape → bitcastoptions.ReshapeIsBitcast(from,to) (0x1dd0e0c0);重塑或复制比特广播链(BitcastingOperandOfReshapeOrCopyChain 0x1dd0df60)
HandleTranspose(底座)身份转置去除(RemoveTranspose 0x...);转置转置组合(SimplifyTranspose 0x...);将转置推向叶子 (TryToReorder… / TryToSink…)
HandleDot0x1dd1d3a0点规范化; OptimizeDot系列(0x1dd...);转置/转换吸收;零/空收缩暗淡折叠;点→减少/乘以强度减少
HandleConvolution0x1dd48280FoldConv…/SwapConv… 操作数折叠;琐碎窗口简化
HandleReduce0x1dd3b540MergeReduces (0x1dd3adc0):相邻熔丝减少了兼容调光;减少广播/减少重塑折叠
HandleConcatenate(底座)单操作数串联消除;相邻常量连续折叠;切片连接重组
HandlePad(底座)零填充消除;垫的垫组成;负填充处理(选项门控)
HandleGetTupleElement(底座)GTE(Tuple(...), i) → operand_i(元组/GTE 折叠)
HandleOptimizationBarrier0x1dd27360一个围绕栅栏重写的处理程序 — 见下文

死元组修剪处理程序(交叉引用)

AlgebraicSimplifierVisitor::HandleOptimizationBarrier (0x1dd27360) 是唯一可以重写“围绕”不透明操作而不跨越它的处理程序。当屏障包裹一个元组时,它会在屏障的用户为 GetTupleElement(操作码 64)时遍历它们,将每个索引标记为活动状态(GTE 具有 >1 个用户,索引条目根,或者产生副作用),丢弃死角,重建较小的 kTuple 操作数,并重新索引幸存的 GTE — 防御性 CHECK(use->opcode() == kGetTupleElement) algebraic_simplifier.cc:5302 守卫走道。这在 optimization-barrier.md(处理程序#6)上进行了端到端记录;它证明了简化器可以对元组腿进行死代码式清理,但永远不会跨栅栏重新排序或合并计算。

TPU 处理程序(17 — 16 个布局感知覆盖 + 1 个仅 TPU)

TPU 子类重写处理程序,其开源主体可以生成具有 TPU 代码生成器无法平铺的布局的张量。每个 TPU 处理程序采用 Target const* 并使用 TpuAlgebraicSimplifierVisitor::IsValidLayout(Shape const&) (0x13443660) 保护重写,并通过 TpuAlgebraicSimplifier::UpdateLayout(Shape*) (0x13444860) 重新标记布局。

TPU 处理器地址TPU 特定行为(已恢复)
HandleDot0x13441ce0点强度降低由 xla_tpu_enable_dot_strength_reduction 门控;仅当生成的形状平铺效果良好时,才使用 ShouldStrengthReduceDotToReduce (0x13443800) 来决定 dot → reduce(multiply)
HandleConvolution0x13442f20转换输入板通过 FoldConvInputPad (0x1343dca0) 折叠 — 将 Pad 吸收到转换的 window 中而不是具体化它
HandleAdd0x1343ddc0布局保留添加规范化; lambda $_1 (0x1343e640) 处理操作数重新排序情况
HandleConcatenate0x1343e8a0concat 标准化,使 concat 在 TPU 有利的(短)轴上保持暗淡
HandleSlice / HandleDynamicSlice / HandleDynamicUpdateSlice0x1343f000 / 0x1343f140 / 0x1343f0c0切片/DUS 简化,重新标记结果布局 (UpdateLayout)
HandleReshape / HandleTranspose / HandleBitcast0x1343f1e0 / 0x1343ffe0 / 0x134401c0重塑/转置/位播折叠在结果的 IsValidLayout 上门控
HandleBroadcast0x13440100广播规范化,实现 TPU 喜欢的目标暗淡(补充单独的 TpuBroadcastRewriter 通道)
HandleCompare / HandleConvert / HandleCopy0x13440540 / 0x13441060 / 0x13441200比较/转换/复制折叠; HandleCopy 删除了开源通道将保留的布局无操作副本
HandlePad0x1343f360TPU 垫简化(与 FoldConvInputPad 配对)
HandleSort0x134414c0排序操作数简化;使用比较器谓词($_1/$_2 lambdas 0x13441a20/0x13441c00)来决定冗余操作数删除
HandleRngBitGenerator0x13443440仅 TPU 处理程序(无基本等效项)- 在 RNG 扩展器系列运行后简化 RNG 位生成器模式

两个共享的 TPU 帮助器支撑这些:SetupDerivedInstruction(HloInstruction*, HloInstruction*, bool) (0x134445c0) 将元数据/布局传播到新创建的替换,而 UpdateLayout (0x13444860) 标记所选的平铺布局,以便重写的操作对于布局后管道立即合法。

注意 — 为什么是 TPU 子类。 简化器在布局分配 (布局分配.md) **之后运行,因此每条指令都已经带有提交的平铺布局。开源的 HandleReshape/HandleTranspose/HandleCopy 可以重写为对于 MXU 的 128×128 几何体布局未定义或可tileable的形状; TPU 覆盖添加了 IsValidLayout 防护和 UpdateLayout 重新标记,因此窥孔重写永远不会破坏 tpu→LLO 合法化器 (tpu-to-llo-ods.md) 所依赖的布局不变性。 40 个非重写处理程序(例如 HandleGetTupleElement、集合和归约系列处理程序)是布局透明的,因此基体按原样健全。


AlgebraicSimplifierOptions — 盖茨

该通道由 xla::AlgebraicSimplifierOptions 结构体配置,该结构体通过引用构造函数传递,并内联复制到 +8 处的通道中。它是开源 XLA 选项结构(与 GPU/CPU 后端使用的结构相同),因此它承载完整的 XLA 门系列;反编译确认了它的存在(dtor 0x1343ba00),HandleReshape咨询的ReshapeIsBitcast(Shape, Shape)访问器(0x1dd0e0c0),以及RunImpl+135读取的run_to_fixed_point bool。与TPU路径相关的门:

选项(语义)它如何控制规则证据
run_to_fixed_pointRunImpl 中启用重新运行到定点循环(否则一次通过/计算)字节位于 +135 处,在 0x13443f40 处读取
is_layout_sensitive使简化器尊重已提交的布局(TPU 在布局后运行它)TPU 覆盖查询 IsValidLayout 0x13443660
enable_dot_strength_reductionHandleDot 中启用 dot → reduce/multiply 重写标志 xla_tpu_enable_dot_strength_reduction(字符串确认); ShouldStrengthReduceDotToReduce 0x13443800
ReshapeIsBitcast(from, to)决定 HandleReshape 是否将重塑折叠为 bitcast访问器 0x1dd0e0c0;从 BitcastingOperandOfReshapeOrCopyChain 0x1dd0df60 调用
enable_conv_simplification大门 HandleConvolution 折叠/交换TPU FoldConvInputPad 0x1343dca0
minmax_propagate_nan控制 MinMaxToClamp 中的 min/max → clamp NaN 语义助手 0x1dd23380
enable_negative_padding_replacement门控 HandlePad 中的负焊盘处理底座 HandlePad本体

注意 - 唯一恢复的特定于 TPU 的 标志字符串xla_tpu_enable_dot_strength_reduction 它在字符串表中逐字显示(例如 "4\n%xla_tpu_enable_dot_strength_reduction" 和裸 xla_tpu_enable_dot_strength_reduction),确认点强度减少是 TPU 门控规则系列。其他行是标准AlgebraicSimplifierOptions布尔成员,其名称为upstream-XLA;它们以结构体字段的形式出现(结构体被构造和复制),但它们的确切字节偏移量和相应的 xla_* 标志字符串单独固定在有限的 grep 预算中。将这些门的“存在”视为已确认(它是上游选项结构),并将每个字段的偏移量视为未恢复。 [置信度:如每行标记。]


未恢复的内容

  • AlgebraicSimplifierOptions 的完整每字段布局。ReshapeIsBitcast(访问器 0x1dd0e0c0)和 run_to_fixed_point(通过 +135)是字节固定的;已知其余布尔门存在(上游结构,由 AddPass 蹦床构造和复制),但未单独跟踪它们的偏移量和标志字符串绑定。 [信心:偏移量较低。]
  • 每个基本处理程序主体。 40 个非重写处理程序运行上游 algebraic_simplifier.cc 逻辑;他们的完整规则算术在这里没有重新导出(它是开源的)。 TPU 覆盖表面和助手是字节锚定的;每个 TPU Handle* 主体的内容都是从帮助程序调用中勾画出来的,而不是彻底反编译。 [信心:对覆盖集和助手的信心很高;每个处理程序重写细节的中等。]
  • 完整的 140 槽 DfsHloVisitorBase 操作码映射。 ~24 个槽从 reloc-addend walk 中标记出来;剩余的槽和确切的纯虚拟集(每个访问者必须实现的操作码,例如HandleParameter/HandleConstant/HandleReduce)需要完整的vtable遍历。 [置信度:已标记槽位的置信度较高,未标记的其余位置的置信度较低。]
  • xla::match:: EDSL 语法。 每个 Handle* 内的匹配表达式都是从 xla::match:: 组合器构建的(二进制范围内有 1,117 个实例化);组合器 ABI 是一个单独的解码器,此处不再复制。

置信度总结

索赔证据
TPU通证是jellyfish::TpuAlgebraicSimplifierAlgebraicSimplifier的子类符号 TpuAlgebraicSimplifier::RunImpl 0x13443f40AddPass<TpuAlgebraicSimplifier,Target const&,AlgebraicSimplifierOptions&> 0x10954400
Dispatch 是操作码键控的虚拟表,不利于模式匹配DfsHloVisitorBase<HloInstruction*> vtable 0x21d2c320,140 个插槽;每个 Handle<Op> 发出的函数
基地访客发出 56 Handle*; TPU 子类发出 17 个(16 个覆盖 + 1 个仅 TPU)56 AlgebraicSimplifierVisitor::Handle* + 17 TpuAlgebraicSimplifierVisitor::Handle* 发出的函数
HandleRngBitGenerator 仅 TPUTPU 0x13443440;没有发出基本 AlgebraicSimplifierVisitor RngBitGenerator 符号
每次计算条目 Run 返回 Changed-boolAlgebraicSimplifierVisitor::Run 0x1dd010e0,签名已恢复
运行到定点循环,上限 50,非致命登录卡住RunImpl 0x13443f40;帽 0x32tpu_algebraic_simplifier.cc:985 日志字符串
单调用者门(跳过大于 1 个调用者的计算)caller_instructions() 尺寸检查(v19 > 1 跳过)在 RunImpl
选项内联存储在通道 +8 处; run_to_fixed_point+135AddPass 蹦床 0x10954400 将选项复制到 +8RunImpl 读取为 +135
TPU 覆盖具有布局感知能力 (IsValidLayout/UpdateLayout)IsValidLayout 0x13443660UpdateLayout 0x13444860SetupDerivedInstruction 0x134445c0
xla_tpu_enable_dot_strength_reduction 门点强度降低字符串表中的标志字符串; ShouldStrengthReduceDotToReduce 0x13443800
源单元 tpu_algebraic_simplifier.cc / …/simplifiers/algebraic_simplifier.cc字符串锚定(第 961/985/991 行;基本文件路径)
完整 AlgebraicSimplifierOptions 字段布局/标志绑定仅 2 个字段字节固定

交叉引用 {#cross-references}

  • HLO 预通过 — 有序预通过集;这是通行证 #60 (TpuAlgebraicSimplifier),位于 HloOptimizeThroughLayoutAssignment 的尾部。
  • 编译器概述 — 五相脊柱;该简化器位于系列 1(HLO 优化过程)的第 1 阶段。
  • 优化屏障 — 文档 HandleOptimizationBarrier(该访问者的处理程序 #6),唯一在不透明操作周围修剪死元组腿的重写。
  • 布局分配 — 提交 TPU 子类的 IsValidLayout/UpdateLayout 所遵守的平铺布局(TPU 子类存在的原因)。
  • HLO 通证注册表 — 此通道是 HloPassInterface RTTI 目录的一个条目。
  • 编译阶段 0–3 — 第 1 阶段 HLO 优化相对于 MLIR 降低交叉点而言。
  • 融合模式 — 简化后运行的融合过程;他们的优先级队列调度程序是一个“第三”引擎,与这个操作码虚函数表访问者不同。
  • 点/转换 → MXU 降低 — 最终降低 Dot/Convolution 操作 TPU HandleDot/HandleConvolution 规范化。
  • 二进制: extracted/libtpu-0.0.40-cp314-cp314-manylinux_2_31_x86_64/libtpu/libtpu.so(构建 ID 89edbbe81c5b328a958fe628a9f2207d)
  • 索引条目: 第五部分 — 编译器:降低和优化通道/前端和管道 — 返回索引