MHLO → XTile → tpu 降低

本页中的所有地址、符号和操作名字符串均适用于来自 libtpu-0.0.40-cp314 wheel 的 libtpu.so（build-id 89edbbe81c5b328a958fe628a9f2207d，构建 libtpu_lts_20260413_b_RC00）。其他版本会有所不同。

摘要

本页记录二进制中实际存在的 MHLO/StableHLO → XTile → tpu 降低，而关于它最重要的事实是结构性的，也是纠偏性的。短语 "MHLO → XTile → tpu" 描述的是两个共享 libtpu.so 但从不相连的不同降低树：

MHLO/StableHLO → XTile 是真实存在的。 XTile (xla::xtile::XTileDialect) 是一个分块张量 MLIR 方言，包含六个注册操作、四个前端降低 pass，以及逐字对应的 StableHLO→arith 模式表。它位于 third_party/.../xla/codegen/xtile/ 下，并通过 xtile-cpu-* 流水线降低到 CPU/LLVM 栈。
XTile → tpu 并不存在。 没有任何 pass 从 XTile 生成 tpu 方言。该二进制包含零个 *ToTpuPass / MhloToTpu / StablehloToTpu 转换 pass 函数。XTile 的依赖方言集合是 CPU/LLVM 集合（xla::cpu::XlaCpuDialect、memref、vector、LLVM），在任何 XTile pass 中都没有出现 tpu、llo 或 sparse_core。

因此，编译器概览中称为 Level 2 的 tpu 方言不是通过降低通用 MHLO 产生的。在 TPU 设备路径上，通用 HLO 由约 3225 个 xla::jellyfish::*Emitter 类直接发射到 LLO；tpu 方言只会被导入：由 Pallas/Mosaic 前端编写，嵌入到 HLO kCustomCall("tpu_custom_call") 中，并由 xla::jellyfish::mlir_utils::GetMlirModuleOpFromCustomCall 提取。本页完整记录树 B（XTile），并明确树 A 的 tpu 方言实际源自何处。认为 "XTile sits between MHLO and tpu" 的解读实际上混淆了两个断开的树：此二进制中的符号、字符串和依赖方言证据都显示不存在 XTile→tpu 链接（见双树图景）。

对于重新实现，本页固定的契约是：

两棵树及其边界。 树 A（TPU 设备：HLO → jellyfish *Emitter → LLO；tpu 仅通过 Mosaic custom-call 导入）对比树 B（XLA CPU/GPU 代码生成：StableHLO/MHLO → XTile → memref/vector → LLVM）。解释为什么 "XTile → tpu" 没有 pass。
四个 XTile 前端 pass：stablehlo-lower-to-xtile、stablehlo-lower-to-arith、convert-elementwise-0d-tensor-to-scalar、xtile-verify-legal-ops；包括它们的 CLI 名称、创建函数、依赖方言，以及在 xtile-cpu-* 流水线中的顺序。
六个注册的 XTile 操作，包括恢复出的 build() 签名，另有两个属性（xtile.layout、xtile.tiling_info）和 TiledBufferInterface。
逐操作的 MHLO/StableHLO → XTile 映射：逐字的 StableHLO→arith 模板表、Emit* fusion-emitter 辅助函数，以及 elementwise、dot、block-scaled dot、transpose、reshape 和 constant 的具体降低示例。
tpu 方言实际来自何处：Mosaic 导入边界（GetMlirModuleOpFromCustomCall / RunMLIRPasses），避免重新实现者把精力浪费在不存在的 MHLO→tpu legalizer 上。


XTile 方言	`xla::xtile::XTileDialect`（`initialize` @ `0x1507ec20`）；源码 `third_party/.../xla/codegen/xtile/ir/xtile_ops.cc`
前端 pass	`StablehloLowerToXtilePass` `0x150602c0` · `StablehloLowerToArithPass` `0x1505a880` · `ConvertElementwise0DTensorToScalarPass` `0x15059440` · `VerifyLegalXTileOpsPass` `0x15062300`
终端 CPU pass	`xla::cpu::LowerXTileEntryPass`（`xtile-cpu-lower-xtile-entry`；create @ `0x14d905c0`）
顶层 emitter	`xla::xtile::EmitXTileModule` @ `0x14c1c9e0`（HLO fusion → `xtile.entry_func` module）
注册操作（6）	`entry_func`、`return`、`extract`、`insert`、`mask`、`dot_scaled`（来自 `XTileDialect::initialize` 的 `addOperations<>`）
属性（2）	`xtile.layout`（`LayoutAttr`）· `xtile.tiling_info`（`TilingInfoAttr`）
XTile 降低目标	CPU/LLVM（`memref`/`vector`/`scf`/`LLVM`），不是 `tpu` 或 `llo`
`tpu` 方言来源	仅导入：`GetMlirModuleOpFromCustomCall` @ `0x13e327a0`；流水线 `RunMLIRPasses` @ `0x111fefa0`
二进制中的 MHLO→`tpu` pass	无（0 个 `*ToTpuPass`/`MhloToTpu`/`StablehloToTpu` 符号）
置信度	高（由符号/字符串锚定），除非某行或标注另有说明

双树图景

支配本页的单一结构性事实是：字符串 "MHLO → XTile → tpu" 命名的是一条并不作为单链存在的链。两个独立的降低树被链接进 libtpu.so，而中间的 tpu 方言链接缺失。

text

TREE A — TPU device (the real product path)
┌────────────────────────────────────────────────────────────────────────┐
│  HLO  ──[97-row pre-passes]──►  HLO                                       │
│       ──[~3225 xla::jellyfish::*Emitter + LloRegionBuilder]──►  LLO       │
│       ──[bundle packer]──►  ISA bundles                                   │
│                                                                          │
│  PLUS, for JAX/Pallas kernels ONLY:                                      │
│  HLO kCustomCall("tpu_custom_call")  carrying a serialized `tpu` module   │
│       ──[GetMlirModuleOpFromCustomCall  0x13e327a0]──►  tpu-dialect IR    │
│       ──[RunMLIRPasses 16-stage pipeline]──►  LLO                         │
└────────────────────────────────────────────────────────────────────────┘
        the `tpu` dialect is authored upstream and imported — never
        produced by lowering MHLO

TREE B — XLA CPU/GPU codegen (bundled, off the TPU device path)
┌────────────────────────────────────────────────────────────────────────┐
│  StableHLO/MHLO  ──[stablehlo-lower-to-xtile]──►  XTile (+arith/math)     │
│       ──[stablehlo-lower-to-arith]──►  XTile + arith/math                 │
│       ──[convert-elementwise-0d-tensor-to-scalar]──►  scalarized          │
│       ──[xtile-verify-legal-ops]──►  (gate)                               │
│       ──[xtile-cpu-bufferization → memref]──►                            │
│       ──[xtile-cpu-*-to-vector / -to-loops]──►  vector/scf/memref         │
│       ──[xtile-cpu-lower-xtile-entry]──►  func/LLVM  ──►  CPU kernel       │
└────────────────────────────────────────────────────────────────────────┘
        XTile lowers to the LLVM/CPU stack, NEVER to `tpu`
```text

所以 "MHLO → XTile" 是树 B，并且完全真实；"XTile → `tpu`" 是断开的链接。证据如下：

| 观察 | 此二进制中的证据 |
|---|---|
| 每个 XTile transform pass 都是 `xtile-cpu-*` | 恢复出的 CLI 字符串：`xtile-cpu-bufferization`、`-fuse-elementwise`、`-linalg-elementwise-to-vector`、`-memref-copy-to-loops`、`-shlo-to-vector`、`-tensor-ops-to-bufferizable`、`-unpack-sub-byte-vector-write`、`-vector-to-scalar`、`-lower-xtile-entry` |
| XTile 位于 XLA *codegen* 树中，而不是 jellyfish | 源路径 `codegen/xtile/ir/xtile_ops.cc`、`codegen/xtile/ir/transforms/lower_stablehlo_to_xtile.cc`、`codegen/xtile/codegen/fusion_emitter.cc`、`…/dot_algorithms.cc`、`…/emitter_helpers.{cc,h}`、`…/tiled_emitter_constraints.cc` |
| XTile 的依赖方言集合是 CPU/LLVM，没有 `tpu` | 注册表 `insert<cf, func, math, xla::cpu::XlaCpuDialect, mhlo, scf, LLVM, tensor, vector, xla::Xla, xla::xtile, stablehlo, linalg, memref, ub>`；确认存在 `XlaCpuDialect`；任何 XTile pass 中都没有 `tpu`/`llo`/`sparse_core` |
| XTile 操作通过 MLIR 标准 One-Shot Bufferize 进行 bufferize | `ExtractTileOp`/`InsertTileOp` 实现 `mlir::bufferization::BufferizableOpInterface`；`tpu`/`llo` 操作从不这样做（它们走 jellyfish 分配） |
| 终端 pass 是 `xla::cpu::` | `LowerXTileEntryPass` 将 `xtile.entry_func` 重写为 CPU/LLVM emitter 使用的 `func.func` |
| 不存在 MHLO→`tpu` pass | grep 函数表：对 `*ToTpuPass`、`*ToTPUDialect`、`MhloToTpu`、`HloToTpu`、`StablehloToTpu`、`LegalizeToTpu`、`LowerMhloToTpu` 的命中数为 **0** |

> **注意： "TLP" 不是方言，也不是 XTile。** [概览](overview.md)中的 "TLP"（TPU-Level Program）是 Phase 2a 的 HLO→MLIR 导入载体的概念名称。不存在 `tlp` MLIR 方言，也不存在 `xla::tlp::`/`mlir::tlp::` 符号。在 TPU 路径上，`tpu`（Mosaic）方言是 "TPU-Level Program" 容器的物理实现，但**仅**用于导入的 Pallas/Mosaic kernel。通用 MHLO 由 jellyfish `*Emitter` 类直接降低到 LLO；它从不会变成 `tpu` 操作。XTile 是一个无关的 CPU/GPU 方言，只是碰巧随同一个 `.so` 发布。

---

## 四个 XTile 前端 Pass

MHLO/StableHLO → XTile 前端由 `xla::xtile::(anonymous namespace)` 中的四个 pass 构成，每个都由 `impl::*PassBase` TableGen 模板生成。前三个重写 tiled fusion 主体；第四个是合法性闸门。所有地址都是 `.text` create-function 入口点。

| Pass 类 | create() | `runOnOperation` | CLI 参数 | 作用 |
|---|---|---|---|---|
| `StablehloLowerToXtilePass` | `0x150602c0` | `0x15060560` | `stablehlo-lower-to-xtile` | 结构性/张量形状的 StableHLO → `xtile.extract`/`insert`/`mask`/`dot_scaled` tile 操作 + `Emit*` tile-dot/transpose/reshape 辅助函数 |
| `StablehloLowerToArithPass` | `0x1505a880` | `0x1505aa60` | `stablehlo-lower-to-arith` | 标量/elementwise StableHLO → `arith`/`math` |
| `ConvertElementwise0DTensorToScalarPass` | `0x15059440` | `0x15059620` | `convert-elementwise-0d-tensor-to-scalar` | `tensor<f32>`（rank-0）→ 标量 `f32` |
| `VerifyLegalXTileOpsPass` | `0x15062300` | `0x150624c0` | `xtile-verify-legal-ops` | 闸门：只允许 XTile/`arith`/`math`/`tensor`/`func` 合法 |

`StablehloLowerToXtilePass` 的描述字符串被逐字恢复为：`"Lowers stablehlo ops to Xtile."`

**依赖方言**（来自每个 `*PassBase::getDependentDialects` 主体）：

- `StablehloLowerToXtilePass`：注册 `arith`、`stablehlo`，再加上共享的 XTile/`tensor`/`math` 注册表插入。**输入** = StableHLO（容忍少量 MHLO-only 操作白名单）；**输出** = XTile + `arith` + `math` + `tensor`。
- `StablehloLowerToArithPass`：注册 `arith`。**输入** = StableHLO 标量/elementwise 操作；**输出** = `arith`/`math`。
- `ConvertElementwise0DTensorToScalarPass`：由 `mlir::TypeConverter` 驱动的*完整转换*，将 rank-0 ranked tensor 映射为标量，带有 source+target materialization 回调，以及 `addDynamicallyLegalOp<arith::ConstantOp>`，用于在 materializer 触发前保持 0-D 常量合法。
- `VerifyLegalXTileOpsPass`：一个遍历模块（`walk<Operation*>`）的 verifier；如果合法 XTile/`arith`/`math`/`tensor`/`func` 集合之外的任何操作存活，就以 `"Could not legalize op: "` 失败。这是 "must be gone before XTile codegen" 闸门。

> **注意：`lower-to-xtile` 和 `lower-to-arith` 会一起运行在 tiled fusion 主体上。** `StablehloLowerToXtilePass` 处理*形状层级*操作（会变成 tile 读/写、tile dot、transpose 和 reshape 的操作），而 `StablehloLowerToArithPass` 处理已经提取出的 tile 上的*标量/elementwise*操作。它们是一次 legalization 的互补两半，而不是各自独立完成全部 legalization 的顺序阶段。

### 在 `xtile-cpu-*` 流水线中的位置

四个前端 pass 是完整 CPU codegen 流水线的第 1–4 阶段，按降低顺序如下（所有 CLI 名称均从字符串池逐字恢复）：

```text
 1  stablehlo-lower-to-xtile                 shape ops → xtile.extract/insert/mask/dot_scaled + Emit* tile-dot/transpose/reshape helpers
 2  stablehlo-lower-to-arith                 scalar/elementwise → arith/math
 3  convert-elementwise-0d-tensor-to-scalar  tensor<f32> → f32
 4  xtile-verify-legal-ops                   gate (only XTile/arith/math/tensor/func legal)
 5  xtile-cpu-fuse-elementwise               fuse adjacent elementwise tile ops
 6  xtile-cpu-shlo-to-vector / -linalg-elementwise-to-vector   → vector dialect
 7  xtile-cpu-tensor-ops-to-bufferizable     normalize stray tensor.* ops
 8  xtile-cpu-bufferization                  tensor → memref (One-Shot Bufferize)
 9  xtile-cpu-memref-copy-to-loops           memref copies → scf.for
10  xtile-cpu-vector-to-scalar               residual vector → scalar
11  xtile-cpu-unpack-sub-byte-vector-write   sub-byte stores
12  xtile-cpu-lower-xtile-entry              xtile.entry_func → func.func (xla::cpu::LowerXTileEntryPass)
13  → LLVM dialect → CPU object code

陷阱：这条流水线不是 97 行 HLO pre-pass 表中的一行。 HLO pre-pass（编译阶段）属于 TPU 设备编译器（DeepseaCompilerBase::RunHloPasses）。xtile-cpu-* 流水线是 XLA CPU 后端的逐 fusion codegen，在 CPU 后端自己的 StableHLO/MHLO 规范化之后由 EmitXTileModule 调用。设备流水线中没有任何一行会添加 xtile-* pass。重新实现者如果在 TPU phase 列表中寻找 XTile 阶段，将找不到；这就是双树分裂的具体后果。

六个注册的 XTile 操作

XTileDialect::initialize（0x1507ec20）通过单个 addOperations<DotScaledOp, EntryFuncOp, EntryFuncReturnOp, ExtractTileOp, InsertTileOp, MaskOp>() 注册六个操作。每个都有 build/verify/print/parse 和一个 xtile.<name> 操作名字符串（xtile.entry_func、.return、.extract、.insert、.mask、.dot_scaled 均从字符串池恢复）。不存在 xtile.dot 操作，见下方陷阱。

操作名	C++ 类	恢复出的 `build()` 签名	作用
`xtile.entry_func`	`xla::xtile::EntryFuncOp`	`(StringRef name, ArrayRef<Type> argTypes, ArrayRef<NamedAttribute>, ArrayRef<DictionaryAttr>)`	tiled-fusion 入口函数（`FunctionOpInterface`）
`xtile.return`	`xla::xtile::EntryFuncReturnOp`	`create(OpBuilder&, Location)`；terminator，无操作数	`entry_func` terminator
`xtile.extract`	`xla::xtile::ExtractTileOp`	`(Type result, Value source, ValueRange offsets, ArrayRef<long> staticOffsets, ArrayRef<long> staticSizes)`	从 tiled buffer/tensor 中读取一个 tile（`BufferizableOpInterface`、`TiledBufferInterface`）
`xtile.insert`	`xla::xtile::InsertTileOp`	`(Value tile, Value dest, ValueRange offsets, ArrayRef<long> staticOffsets, ArrayRef<long> staticSizes)`	将 tile 写回 tiled buffer/tensor（`BufferizableOpInterface`、`TiledBufferInterface`）
`xtile.mask`	`xla::xtile::MaskOp`	`(Value source, ArrayRef<long> maskedDims, Value padValue)`	对部分边界 tile 做谓词化/填充；有 `fold()` + `inferReturnTypes`
`xtile.dot_scaled`	`xla::xtile::DotScaledOp`	`(Type, Value lhs, Value rhs, Value lhsScale, Value rhsScale, bool, bool, bool)`	block-scaled（MXFP）tile dot；自定义 print/parse

注意：xtile.extract/xtile.insert 使用规范的混合静态/动态 slice 形式。 它们的操作数是 static-offset + dynamic-offset（ValueRange）+ static-size，形状与 tensor.extract_slice/tensor.insert_slice 完全一致。这正是它们能在 One-Shot Bufferize 下干净地 bufferize 为 memref.subview 风格读取和原地 tile store 的原因。
陷阱：不存在普通的 xtile.dot 操作。 addOperations<> 只注册上面的六个操作；二进制中没有 xla::xtile::DotOp 类、没有 xtile.dot 操作名字符串，也没有 DotOp::build/print/parse 符号。普通（非 scaled）的 tile dot 不是 XTile 操作，而是由 dot-algorithm emitter EmitSingleTileDot（0x14c277a0，dot_algorithms.cc）直接发射到 linalg/vector。只有 block-scaled DotScaledOp 是一等操作。逐算法的 dot 指令序列（f32×f32、bf16×bf16→f32、tf32、x3/x6 高精度）位于 dot_algorithms.cc，此处不枚举；对精确逐算法主体的置信度为低。

XTile 属性和接口

符号	编码内容
`xla::xtile::LayoutAttr`（`xtile.layout`）	minor-to-major 布局：`minor_to_major : DenseI64ArrayAttr`（解析错误字符串 `"failed to parse XTile_LayoutAttr parameter 'minor_to_major' …"`）
`xla::xtile::TilingInfoAttr`（`xtile.tiling_info`）	`tile_count : int32_t`、`tiles_per_workgroup : int32_t`
`xla::xtile::TileInfo`	由 `EmitParameterExtract` 消费的逐操作 tile 描述符
`xla::xtile::TiledBufferInterface`	由 `ExtractTileOp` + `InsertTileOp` 实现的操作接口
`xla::xtile::DotOperands` / `ScaledDotOperands` / `DotOperandSide`	dot 操作数路由的辅助结构

tiles_per_workgroup 字段加上 CPU-thunk workgroup 字符串（NumWorkGroups{%d, %d, %d}、XLA_CPU_NumWorkGroups、xla.cpu.KernelThunkProto）确认了 CPU/GPU workgroup 并行模型：某个维度被拆成 tile_count 个 tile，每个并行 workgroup 分配 tiles_per_workgroup 个 tile。这是 CPU/GPU codegen 概念；TPU 侧使用的是 sublane/lane vreg tiling（Mosaic 布局推断）。

逐操作 MHLO/StableHLO → XTile 映射

XTile 降低有清晰分工：标量/elementwise 操作走逐字的 StableHLO→arith 表（StablehloLowerToArithPass），而形状层级操作走 Emit* fusion-emitter 辅助函数（StablehloLowerToXtilePass、fusion_emitter.cc/emitter_helpers.cc）。

StableHLO → arith 表（逐字）

每一行都是从 demangled 名称 LowerStableHloOpToArith<Src, FloatDst, SignedIntDst[, UnsignedIntDst]> 恢复出的模板实例。Float 与 int 派发在匹配时由操作数元素类型选择。

StableHLO 源	float 目标	signed-int 目标	unsigned-int 目标
`stablehlo.add`	`arith.addf`	`arith.addi`	`arith.addi`
`stablehlo.subtract`	`arith.subf`	`arith.subi`	`arith.subi`
`stablehlo.mul`	`arith.mulf`	`arith.muli`	`arith.muli`
`stablehlo.div`	`arith.divf`	`arith.divsi`	`arith.divui`
`stablehlo.rem`	`arith.remf`	`arith.remsi`	`arith.remui`
`stablehlo.max`	`arith.maximumf`	`arith.maxsi`	`arith.maxui`
`stablehlo.min`	`arith.minimumf`	`arith.minsi`	`arith.minui`
`stablehlo.and`	`arith.andi`	`arith.andi`	`arith.andi`
`stablehlo.or`	`arith.ori`	`arith.ori`	`arith.ori`
`stablehlo.xor`	`arith.xori`	`arith.xori`	`arith.xori`

一元操作通过 LowerStableHloUnaryOpToMath<Src, MathDst> 路由（恢复出的实例：stablehlo.round_nearest_even → math.roundeven）。同一匿名 namespace 中的额外专用模式类：

LowerCompareOp：stablehlo.compare → arith.cmpf/arith.cmpi（谓词由比较方向 + 类型决定）。
LowerConvertOp：stablehlo.convert → arith.ext*/trunc*/sitofp/uitofp/fptosi/fptoui（类型对表；类已确认，完整逐类型对映射为推断，置信度中等）。
ConstantConversionPattern：stablehlo.constant → arith.constant（DenseElementsAttr 透传）。
ElementwiseConverter：通用 elementwise dispatcher。

形状层级 emit 辅助函数

这些函数（位于 xla::xtile:: / 其匿名 namespace）构建由 arith 操作填充的 tile-op 结构：

辅助函数	text 地址	构建内容
`EmitXTileModule`	`0x14c1c9e0`	顶层：fusion 的整个 `xtile.entry_func` module
`EmitScope`	`0x15066ec0`	遍历 HLO instruction span，将 XTile/`arith` 发射到 value map
`EmitParameterExtract`	`0x15066e00`	根据 `TileInfo` 对参数 tile 执行 `xtile.extract`
`EmitConstant`	`0x15064da0`	为 HLO 常量发射 `arith.constant`/splat
`EmitElementwise`	`0x150630a0`	elementwise HLO → 提取出的 tile 上的 `arith`/`math`
`EmitSingleTileDot`	`0x14c277a0`	`dot` 的一个 tile → dot 算法（`dot_algorithms.cc`）
`EmitSingleTileScaledDot`	`0x14c289c0`	block-scaled dot 的一个 tile → `xtile.dot_scaled`
`EmitTiledReshape`	`0x150694c0`	tiled tensor 上的 reshape
`EmitTiledTranspose`	`0x15069860`	tiled tensor 上的 transpose（`SmallVector<long,6>` permutation）
`EmitTiledComputation` / `EmitTiledInstructionList` / `EmitTiledBitcast` / `EmitReduceComputation` / `EmitNestedFusion`	（多个符号）	整个 computation / 逐 instruction / bitcast / reduce body / nested-fusion 内联

EmitXTileModule 签名（已恢复）：

cpp

EmitXTileModule(std::string_view name,
                HloFusionInstruction const*,
                SymbolicTileAnalysis const&,
                Tiling const&,
                MLIRContext&,
                absl::Span<Type>,
                std::optional<stream_executor::GpuComputeCapability> const&);
```text

XTile module 使用 symbolic tile analysis + 选定 tiling，直接从 HLO **fusion** 发射，并受 GPU compute capability 约束（CPU 传入 `nullopt`）。入口日志锚点：`"Emitting XTile IR for fusion"`。Dot emit 规范化使用 `CanonicalizeDotOperand`（`0x150692a0`）、`MaskDotOperand`、`GetDotLoopIterationCount`、`GetDotAccumulatorType`、`GetPaddedTileSizes`。

### 具体降低示例

```mlir
// (a) elementwise binary — stablehlo.add of two tiled f32 tensors
//     %c = stablehlo.add %a, %b : tensor<256x256xf32>
//     after tiling + StablehloLowerToArith over one tile:
%ta = xtile.extract %A[%i, %j] [/*static offsets*/] [/*static sizes*/] : tensor<TMxTNxf32>
%tb = xtile.extract %B[%i, %j] [...] [...]                              : tensor<TMxTNxf32>
%tc = arith.addf %ta, %tb : tensor<TMxTNxf32>      // float vs int per the arith table
      xtile.insert %tc into %C[%i, %j] [...] [...]

mlir

// (b) dot (matmul), one tile — EmitSingleTileDot (dot_algorithms.cc)
//     stablehlo.dot_general %lhs, %rhs {contracting = ...}
//     per output tile: extract lhs/rhs tiles; loop GetDotLoopIterationCount over the
//     contraction tiling; accumulate in GetDotAccumulatorType (f32 acc for bf16/f8 inputs);
//     operands canonicalized by CanonicalizeDotOperand, boundary-masked by MaskDotOperand;
//     result xtile.insert into the output buffer.
```text

```mlir
// (c) block-scaled dot (MXFP) — EmitSingleTileScaledDot
//     a block_scaled_dot custom-call (scale type f8E8M0FNU):
%d = xtile.dot_scaled %lhsTile, %rhsTile, %lhsScale, %rhsScale
       {transpose_lhs, transpose_rhs, acc}    // three bool flags
// verifier strings: "expect scale operands dimension 2 to equal C/block_size",
// "block sizes for cast_from_block_scaled and cast_to_block_scaled must match"

mlir

// (d) transpose — EmitTiledTranspose(b, permutation, tiledValue)
//     stablehlo.transpose {permutation} → tile-level transpose over the tiled
//     RankedTensorType using the SmallVector<long,6> permutation.

// (e) reshape — EmitTiledReshape(b, newShape, tiledValue)
//     stablehlo.reshape → tile-level reshape; collapse/expand of unit dims via CollapseUnitDims.

// (f) constant — EmitConstant → arith.constant (DenseElementsAttr); 0-D constants are
//     later scalarized by ConvertElementwise0DTensorToScalarPass.
```text

> **注意：broadcast / iota / concatenate。** `BroadcastInDims`、`Splat` 和 `ConstLike` 辅助函数处理 broadcast；iota 和 concatenate 走 `EmitGeneric` / `CheckConcatenateOperands`。上游 StableHLO→Linalg 的 `IotaConverter`/`ConcatenateConverter` 类也随二进制发布，但它们属于 GPU/CPU `stablehlo-to-linalg` 路径，**不是** XTile pass；不要混淆。

---

## Tiling 和类型转换

### XTile 消费 tiling；它不选择 tiling

XTile **不**决定 tiling。这个决定更早由 `xla::SymbolicTileAnalysis` 做出：

- `AnalyzeFusion`（`0x14c35a40`）/ `AnalyzeComputation`（`0x14c359a0`）在 fusion 中传播 `IndexingMap`。
- `GetValidTilings`（`0x14c40280`）枚举有效 `Tiling`；`ParametersSatisfyConstraints`（`0x14c3d040`）+ `tiled_emitter_constraints.cc` 对它们进行过滤。
- `ComputeOutputTilingInfo`（`0x14c46960`）/ `ComputeTiledInstructions`（`0x14c47f80`）生成逐 instruction 的 `TileInfo`。

选定的 `Tiling` + `GpuComputeCapability` 会传入 `EmitXTileModule`，后者把 `TilingInfoAttr{tile_count, tiles_per_workgroup}` 和 `LayoutAttr{minor_to_major}` 烘进入口函数和 tile。失败锚点：`"cannot compute parametric tile sizes for dynamically shaped payload op"`。这是 GPU/CPU symbolic-tiling 搜索；TPU 类比物（window-config 搜索、Mosaic sublane/lane tiling）是完全独立的代码。

### 两种类型转换

1. **`tensor<f32>` → 标量 `f32`**（`ConvertElementwise0DTensorToScalarPass`）：一个 `mlir::TypeConverter` 完整转换，包含 `addConversion(Type)`（identity）、`addConversion(RankedTensorType)`（rank-0 → element type）、`addSourceMaterialization`（scalar → tensor re-wrap）、`addTargetMaterialization`（tensor → scalar unwrap），以及 `addDynamicallyLegalOp<arith::ConstantOp>`，让 0-D 常量在 materializer 触发前保持合法。

2. **`tensor` → `memref`**（`xtile-cpu-bufferization`）：通过 `ExtractTileOp`/`InsertTileOp` 实现的 `BufferizableOpInterface` 执行 MLIR One-Shot Bufferize：`ExtractTileOp::bufferizesToMemoryRead = true` 作用于 source，并发射 `memref.subview` 风格读取；`InsertTileOp::bufferizesToMemoryWrite = true` 作用于 dest，并发射原地 tile store；`bufferizesToAllocation` 控制 scratch 分配。`xtile-cpu-tensor-ops-to-bufferizable` 先规范化残留的 `tensor.*` 操作；`xtile-cpu-memref-copy-to-loops` 随后将 copy 降低到 `scf` loop。入口函数 memref 必须使用 tiled layout；锚点：`"All memref arguments should use the TiledLayoutAttr for layout"`；对齐检查使用 `"Cannot confirm that the memref is memory tile-aligned …"`。

---

## XTile 边界处失败/不支持的操作

`xtile-verify-legal-ops` 和两个 lower-to-* 完整转换会拒绝合法集合之外的任何内容。恢复出的失败锚点：

| 锚点字符串 | 触发条件 |
|---|---|
| `"Could not legalize op: "` | 没有 lower-to-xtile/arith 模式的操作存活 |
| `"cannot compute parametric tile sizes for dynamically shaped payload op"` | 动态形状 payload 没有静态 tiling |
| `"Boundaries of the clamp are not legal: "` | `xtile.mask`/clamp 边界越界 |
| `"All tiles must have the same rank."` | 混合 rank 的 tile 输入 tile op |
| `"At least one tile shape must be specified."` | tiling 信息缺失 |
| `"All tiled squeezed dimensions must be of size 1."` | tiled 维度上的 squeeze 非法 |
| `"All memref arguments should use the TiledLayoutAttr for layout"` | entry-func memref 缺少 tiled layout |
| `"Cannot confirm that the memref is memory tile-aligned …"` | 对齐检查无法证明 producer 是 tile-aligned |
| `"expect scale operands dimension 2 to equal C/block_size "` | `xtile.dot_scaled` scale 操作数形状不匹配 |
| `"Allow XLA's MHLO ops not in StableHLO to remain present after legalization (copy, add_dependency, fusion, etc.)"` | 标志：容忍 MHLO-only 操作；若未降低，之后会命中 `"Could not legalize op:"` |

"MHLO not in StableHLO" 标志确认输入是由 MHLO 生成的 StableHLO，并允许一小组 MHLO-only 操作（`copy`、`add_dependency`、`fusion`）持续到 XTile 降低中。

---

## `tpu` 方言实际来自何处

因为不存在 MHLO→`tpu` pass，重新实现者需要知道 TPU 路径上 `tpu` 方言的*真实*来源：它是导入的，不是降低生成的。边界如下：

1. 一个 Pallas `@pl.kernel`（或手写 Mosaic kernel）由 **JAX 前端在 `libtpu` 外部**编译成序列化的 `tpu` 方言 MLIR module。
2. 该 module 被嵌入到 `HloOpcode::kCustomCall` 中，其中 `custom_call_target = "tpu_custom_call"`，位于 backend config 的 `mlir_module` 字段。
3. 在 HLO→LLO 发射时，`xla::jellyfish::mlir_utils::GetMlirModuleOpFromCustomCall`（`0x13e327a0`）提取并缓存 `mlir::ModuleOp`（以 CityHash128 为键，每个唯一 kernel 解析一次）。
4. 导入的 module 运行 `tpu` pass 流水线 `RunMLIRPasses`（`0x111fefa0`）：`MosaicSerdePass`（版本升级）→ memref/vector-layout pass → `createLowerToLLOPass`（[tpu → LLO](tpu-to-llo-ods.md)）→ LLO。

因此，`tpu` 方言是**由上游编写并导入的**，从不由通用 MHLO 降低生成。有关导入/serde/流水线细节，见 [Mosaic 概览](mosaic-overview.md)；有关方言本身，见 [tpu MLIR 方言](tpu-dialect-and-ops.md)。

> **陷阱：不要构建 MHLO→`tpu` legalizer；它不存在。** [概览](overview.md)在 Level 2（`tpu` 方言）绘制的汇合点实际只有两个生产者：Mosaic custom-call 导入（树 A2），以及在通用路径上什么都没有，因为通用 HLO 完全绕过 `tpu`，由 jellyfish `*Emitter` 类直接发射到 LLO（树 A1）。XTile（树 B）两者都不供给；它的目标是 CPU/LLVM。把精力投入 MHLO/StableHLO→`tpu` 转换模式集的重新实现，构建的是生产编译器并不包含的阶段。

---

## 置信度摘要

| 断言 | 证据 |
|---|---|
| 四个 XTile 前端 pass 及所列 CLI 名称/create 函数 | `StablehloLowerToXtilePass` `0x150602c0`、`StablehloLowerToArithPass` `0x1505a880`、`ConvertElementwise0DTensorToScalarPass` `0x15059440`、`VerifyLegalXTileOpsPass` `0x15062300`；CLI 字符串已恢复 |
| 六个注册的 XTile 操作及所列类 | `XTileDialect::initialize` 中的 `addOperations<DotScaledOp, EntryFuncOp, EntryFuncReturnOp, ExtractTileOp, InsertTileOp, MaskOp>`；操作名字符串 `xtile.entry_func/return/extract/insert/mask/dot_scaled` |
| StableHLO→`arith` 表（10 行二元操作） | 从 demangled 名称恢复出的 `LowerStableHloOpToArith<…>` 模板实例 |
| `Emit*` fusion-emitter 辅助函数映射 + `EmitXTileModule` 签名 | `EmitXTileModule` `0x14c1c9e0`、`EmitSingleTileDot` `0x14c277a0`、`EmitTiledTranspose` `0x15069860` 等；日志锚点 "Emitting XTile IR for fusion" |
| XTile 属性 `xtile.layout`/`xtile.tiling_info` 和 workgroup 模型 | `XTile_LayoutAttr`/`minor_to_major` 及 `XTile_TilingInfoAttr`/`tile_count`/`tiles_per_workgroup` 解析错误字符串；CPU-thunk `NumWorkGroups{%d, %d, %d}`/`XLA_CPU_NumWorkGroups`/`xla.cpu.KernelThunkProto` 字符串 |
| `tensor`→`memref` 通过 `BufferizableOpInterface` 上的 One-Shot Bufferize | `ExtractTileOp`/`InsertTileOp` 实现 `BufferizableOpInterface` + `TiledBufferInterface`；`xtile-cpu-bufferization` CLI 字符串 |
| XTile 边界处的失败锚点 | 10 个锚点字符串均从反编译输出恢复 |
| **XTile 不在 TPU MXU 路径上** | XTile 依赖集合 = `XlaCpuDialect`/`memref`/`vector`/`LLVM`（无 `tpu`/`llo`）；`xtile-cpu-*` 流水线；`codegen/xtile/` 源路径 |
| **不存在 MHLO/HLO→`tpu` 转换 pass** | grep 函数表：对 `*ToTpuPass`/`MhloToTpu`/`StablehloToTpu`/`LegalizeToTpu` 的命中数为 0 |
| `tpu` 方言仅通过 `tpu_custom_call` 导入 | `GetMlirModuleOpFromCustomCall` `0x13e327a0`、`RunMLIRPasses` `0x111fefa0`、`MosaicSerdePass`、`tpu_custom_call` 均存在 |
| 普通 tile dot 发射到 `linalg`/`vector`，而不是 `xtile.dot` 操作 | 二进制中没有 `DotOp` 类 / `xtile.dot` 字符串；`EmitSingleTileDot`/`dot_algorithms.cc` 主体未逐算法反编译 |
| 完整 `LowerConvertOp` ext/trunc/itofp/fptoi 类型对表 | `LowerConvertOp` 类存在（119 个引用）；逐类型对映射为推断，未枚举 |

---

## 交叉引用

- [TPU 编译器（概览）](overview.md)：五阶段主干和 IR 层级栈；本页将其中 Level 2（"XTile between MHLO and `tpu`"）细化为双树分裂。
- [编译阶段 0–3](compile-phases.md)：设备编译器的 phase 流水线，`xtile-cpu-*` 流水线**没有**接入其中。
- [tpu MLIR 方言](tpu-dialect-and-ops.md)：XTile **不会**产生的 `tpu` 目标方言；TPU 侧 optimizer 和 Mosaic 路径的汇合点。
- [tpu → LLO 降低](tpu-to-llo-ods.md)：*导入的* `tpu` module 会经过的 `createLowerToLLOPass` 下降路径（真实路径，不同于 XTile 的 CPU/LLVM 目标）。
- [DialectConversion Legalizer](dialect-conversion-legalizer.md)：支撑 XTile 完整转换和 `tpu` legalizer 的 MLIR dialect-conversion 机制。
- [Mosaic 概览](mosaic-overview.md)：Pallas/Mosaic 导入边界（`GetMlirModuleOpFromCustomCall` / `RunMLIRPasses`），也就是 `tpu` 方言的实际来源。
- [Mosaic 布局推断](mosaic-layout-inference.md)：TPU sublane/lane vreg tiling，是 XTile CPU/GPU workgroup tiling 的类比物。
- **二进制：** `extracted/libtpu-0.0.40-cp314-cp314-manylinux_2_31_x86_64/libtpu/libtpu.so`（build-id `89edbbe81c5b328a958fe628a9f2207d`）
- **索引项：** Part V — Compiler: Lowering & Optimization Passes / MLIR lowering chain — [返回索引](../index.md)

MHLO → XTile → tpu 降低 ​

摘要 ​

双树图景 ​

六个注册的 XTile 操作 ​

XTile 属性和接口 ​

逐操作 MHLO/StableHLO → XTile 映射 ​

StableHLO → arith 表（逐字） ​

形状层级 emit 辅助函数 ​