tpu MLIR Dialect：Ops 与 Op-Model 契约
本页中的所有地址、符号名和计数都适用于
libtpu-0.0.40-cp314wheel 中的libtpu.so（build-id89edbbe81c5b328a958fe628a9f2207d，buildlibtpu_lts_20260413_b_RC00）。.symtab未被 strip，因此下面的每个断言都锚定到一个 demangled 符号或 relocation addend。其他版本会有所不同；尤其是 23-slot 顺序固定于这个二进制的 LLVM SHA。
摘要
tpu dialect 是 TPU 的机器 dialect，也是优化器路径（HLO → MHLO → tpu）和 Mosaic 旁路通道（tpu_custom_call → tpu）在下降到 LLO 之前共同汇合的 MLIR 目标 dialect（见编译器概览）。本页从 stripped 但仍带符号的二进制中仍能保留下来的角度记录该 dialect：不是源代码层面的 ODS .td 文件（wheel 中没有这些文件），而是每个 MLIR dialect 中每个 op 都会被编译成的运行时 op-registration ABI。
具体来说，MLIR 注册一个 op 并不是靠 C++ 继承，而是为每个 op 实例化一个类型擦除的 concept 类 mlir::RegisteredOperationName::Model<Op>，并把它作为该 op 的 OperationName::Impl 插入。每个 Model<Op> 都是一个带固定 23-slot 契约的 vtable：一个共享基类析构函数、一个逐 op 的 deleting destructor，以及 21 个分派 slot（fold、canonicalize、verify、parse/print、四个 inherent-attribute accessor，以及九个 property-management 方法）。该二进制包含 6,050 个这样的 Model<Op> vtable，即 60 多个 dialect 命名空间中每个已注册 op 各一个；其中 tpu dialect 拥有 86 个。
全二进制范围的 Model<Op> ABI 由相邻的 MLIR Op-Model 契约 页面负责；本页是它在 tpu dialect 上的具体实例，只在读取 86 个 tpu op 所需范围内复现 23-slot 遍历。本页负责三件事：
Model<Op>23-slot 契约：每个 slot 都标注到其 demangledModel<Op>::<method>符号，说明 1 个共享 + 22 个逐 op 的布局，以及每个 slot 的作用。（已在Model<mlir::tpu::IotaOp>上验证：全部 21 个 dispatch-slot 符号均存在。）- Op↔Model 绑定：一个具体
tpuop 如何填充它的 23 个 slot：Dialect::addOperations<…>registrar、单个RegisteredOperationName::insert，以及从类型擦除 Model thunk 到该 op 的 ODS 生成 static 的三级委托。 tpuop-family 分类：按其OpInterface签名分组的 86 个 op，fold/canonicalize 统计（该 dialect 全部 in-dialect rewrite 表面只有 15 个方法），以及携带可变参数AttrSizedOperandSegments大小数组的 9 个 op（它们属于 53 个Properties结构体非平凡 op 的子集）。
| Op-registration concept | mlir::RegisteredOperationName::Model<Op>（每个已注册 op 一个） |
| Model vtable 计数（所有 dialect） | 6,050，跨 60 多个命名空间（= 已注册 op 总数） |
tpu dialect Model 计数 | 86（精确的 addOperations<> arity；修正了朴素 *Op 字符串扫描得出的约 157 的过计数，见下文） |
| 每个 Model 的 vtable slot | 23 = slot 0 共享 base dtor + slot 1 逐 op dtor + slot 2–22 dispatch |
| 共享 slot-0 dtor | mlir::OperationName::Impl::~Impl（addend 0xfea8820，在全部 6,050 个 Model 中相同） |
tpu op registrar | Dialect::addOperations<mlir::tpu::AllReduceOp, …> @ 0x14aa2c40（86 个实参） |
| 注册 sink | RegisteredOperationName::insert(unique_ptr<OperationName::Impl>, ArrayRef<StringRef>) @ 0x1d8c57a0 |
| Dialect ctor | mlir::tpu::TPUDialect::TPUDialect(MLIRContext*) @ 0x14a96d40（注册 "tpu"、dialect-id 3，然后内联调用 addAttributes<13>/addType<3>/addOperations<86>，没有单独的 initialize()） |
tpu MemoryEffect interface fan-out | 59 个 op（用两种方式交叉核对） |
tpu in-dialect rewrite 表面 | 6 个真实 fold() + 9 个真实 getCanonicalizationPatterns() = 15 个方法 |
带非平凡 Properties 的 tpu op | 53 个（slot 14 getOpPropertyByteSize > 0）；其中 9 个携带 AttrSizedOperandSegments 大小数组 |
| 置信度 | CONFIRMED（字节锚定），除非某一行或标注另有说明 |
Model<Op> 23-Slot 契约
MLIR 的 OperationName 是指向 interned、逐 op OperationName::Impl 的句柄。对于一个已注册 op（即 dialect 声明的 op，而不是从文本解析出的 opaque/unregistered op），该 Impl 是一个 RegisteredOperationName::Model<Op>，也就是每个 op 类型实例化一次的 C++ 模板。Model<Op> 是典型的类型擦除 concept：它派生自抽象的 OperationName::Impl 接口，并用知道具体 Op 的 ODS 生成 static 的 thunk 覆盖每个 virtual。每个已注册 op 恰好有一个 Model<Op> 对象，因此也恰好有一个 vtable；nm 统计到 6,050 个这样的对象（_ZTVN4mlir23RegisteredOperationName5ModelI…EEE）。
该 vtable 遵循 Itanium PIE 布局：vtable+0x00 是 offset-to-top（0），vtable+0x08 指向 _ZTI…Model<Op>… typeinfo，而 vtable+0x10 是 address point，即对象 vptr 存储的值，也是 slot 0 的位置。Slot i 位于 vtable+0x10+8i。在磁盘镜像中，slot word 为零；加载时它们由 R_X86_64_RELATIVE relocation 填充，其 addend 是方法地址。把每个 addend 对照（排序后的）符号表解析，就能恢复每个 slot 的方法名。
下面的参考遍历是 Model<xla::PureCallOp> @ 0x219d4e38；它在排序和签名上与 Model<mlir::tpu::IotaOp> @ 0x219abb28、Model<mlir::llo::ConstantOp> @ 0x2193dba0 和 Model<mlir::mosaic_sc::RelayoutOp> @ 0x21905130 字节级一致，即三个 dialect、一个 ABI。
RegisteredOperationName::Model<Op> vtable (address point = +0x10 = slot 0)
slot off method (Model<Op>::…) nature
---- ---- ------------------------------------------ --------------------------------
0 0x00 OperationName::Impl::~Impl SHARED base dtor (0xfea8820 ∀ ops)
1 0x08 ~Model [deleting] per-op dtor
2 0x10 foldHook(Operation*, ArrayRef<Attribute>, per-op; wraps Op<>::getFoldHookFn
SmallVectorImpl<OpFoldResult>&) lambda (UniqueFunction)
3 0x18 getCanonicalizationPatterns( per-op; no-op unless op defines it
RewritePatternSet&, MLIRContext*)
4 0x20 hasTrait(TypeID) trait / interface membership test
5 0x28 getParseAssemblyFn() returns the op's custom parser fn
6 0x30 populateDefaultAttrs(OperationName const&, default-attr init
NamedAttrList&)
7 0x38 printAssembly(Operation*, OpAsmPrinter&, per-op; wraps getPrintAssemblyFn
StringRef)
8 0x40 verifyInvariants(Operation*) wraps Op<>::verifyInvariants
9 0x48 verifyRegionInvariants(Operation*) region-structure verify
10 0x50 getInherentAttr(Operation*, StringRef) → Op::getInherentAttr (TAIL-JMP)
11 0x58 setInherentAttr(Operation*, StringAttr, → Op::setInherentAttr
Attribute)
12 0x60 populateInherentAttrs(Operation*, → Op::populateInherentAttrs
NamedAttrList&)
13 0x68 verifyInherentAttrs(OperationName, inherent-attr verify
NamedAttrList&, function_ref<…>)
14 0x70 getOpPropertyByteSize() INLINED sizeof(Op::Properties)
15 0x78 initProperties(OperationName, PropertyRef, placement-construct props
PropertyRef)
16 0x80 deleteProperties(PropertyRef) destruct props
17 0x88 populateDefaultProperties(OperationName, default props
PropertyRef)
18 0x90 setPropertiesFromAttr(OperationName, attr → props
PropertyRef, Attribute, fn_ref<…>)
19 0x98 getPropertiesAsAttr(Operation*) props → attr (print / bytecode)
20 0xa0 copyProperties(PropertyRef, PropertyRef) clone props
21 0xa8 compareProperties(PropertyRef,PropertyRef) props equality (for CSE)
22 0xb0 hashProperties(PropertyRef) props hash (for CSE / dedup)
```text
这 23 个 slot 聚成四个功能组：
- **Slot 2–9 — 分发表面。** Folding、canonicalization、trait/interface gate（`hasTrait`）、自定义 parse/print、默认 attribute 填充，以及两个 verifier。Slot 2、3、7 和 8 是*可选* hook；它们对每个 op 都存在，但除非该 op 确实定义了相应方法，否则会包装一个 no-op（见绑定小节）。
- **Slot 10–13 — inherent-attribute accessor。** Inherent attribute 是该 op 的*声明式* attribute（例如 `iota` 的 `dimensions`），区别于 discardable attribute。这四个 slot 直接 tail-jump 到该 op 的 ODS static。
- **Slot 14–22 — property 管理。** MLIR 的 `Properties` 是一个 op 的 inherent attribute 和可变操作数 segment size 的内联存储。这九个 slot 构造、析构、复制、比较、哈希，并对该存储进行（反）序列化。对于没有 `Properties` 结构体的 op，它们是平凡的（见下面的 slot 14）。
> **说明 — “23-slot Model” 在结构上是 1 个共享 + 22 个逐 op，而不是 23 个逐 op。** Slot 0，即每个对象 vptr 存储的 address point，在全部 6,050 个 Model 中都是*同一个* relocation addend `0xfea8820`（`mlir::OperationName::Impl::~Impl`）；这一点已在横跨 `tpu`、`llo`、`sparse_core`、`mosaic_sc` 和 `xla` 的 30-Model 分层样本上验证。它是公共 `OperationName::Impl` 基类子对象的析构函数。只有 slot 1–22 是逐 op 的模板实例化：slot 1 是 `Model<Op>::~Model`（deleting dtor），slot 2–22 是 dispatch thunk。重新实现者应把 `Model<Op>` 建模为*一个共享基类析构函数加 22 个 op-specific 条目*。
>
> **怪异点 — `getOpPropertyByteSize`（slot 14）是内联的 `sizeof(Op::Properties)` 常量，而不是函数调用。** 编译器把该 slot 折叠成一个字面量。`Model<mlir::tpu::IotaOp>::getOpPropertyByteSize` 是 `mov $0x8; ret`，即 8 字节，单个 `dimensions` `DenseArrayAttr` 句柄。`Model<mlir::tpu::BarrierOp>::getOpPropertyByteSize`（@ `0x14ab1960`）是 `xor %eax,%eax; ret`，即 0 字节，因为 `BarrierOp` 没有 `Properties`。这是判断“这个 op 是否有内联 property”的最廉价诊断方式：读取 slot 14 的两条指令函数体。九个 property-management slot（14–22）只有在 op 携带 inherent attribute 或 `AttrSizedOperandSegments` 时才做真实工作。
---
## Op↔Model 绑定
一个具体 op 不是通过在手写子类中覆盖 virtual 来绑定到它的 23-slot Model，而是通过**模板实例化 + 委托**，分三层完成。
### 第 1 层 — 注册：`addOperations` → `insert`
每个 dialect 都用单个可变参数 `Dialect::addOperations<Op…>` 模板实例化来注册其 op。对于 `tpu`，它直接在 dialect 构造函数体中调用（`TPUDialect::TPUDialect` @ `0x14a96d40`，没有单独的 `initialize()` 符号；构造函数运行 `Dialect::Dialect(this, "tpu", 3, …)`，随后依次执行 `addAttributes<…>`、`addType<…>` ×3、`addOperations<…>`）：
```text
mlir::Dialect::addOperations<mlir::tpu::AllReduceOp, mlir::tpu::AllocaSemaphoreOp,
mlir::tpu::AssumeLayoutOp, …> @ 0x14aa2c40它恰好有 86 个类型实参，这确认了 tpu dialect 的 op 数来自 registrar arity，而不是字符串扫描。对于 pack 中的每个 Op，addOperations 会 materialize 一个 Model<Op>（其实例的 vtable 持有 23 个 Model<Op>::* thunk），并把它交给单一、非模板的注册 sink：
mlir::RegisteredOperationName::insert(
std::unique_ptr<mlir::OperationName::Impl, …>, ← the Model IS the Impl
llvm::ArrayRef<llvm::StringRef>) @ 0x1d8c57a0
```text
`Model<Op>` *就是* `OperationName::Impl`，注册的含义是“把这个 Impl 以该 op 的名称 intern 起来”。`TPUDialect` 构造函数（`mlir::tpu::TPUDialect::TPUDialect(MLIRContext*)` @ `0x14a96d40`）在自身函数体中把 attributes（`addAttributes<13>`）、types（`addType<…>` ×3）和 ops（`addOperations<86>`）全部接好；`initialize()` 完全内联进构造函数，因此没有独立的 `TPUDialect::initialize` 符号。全二进制范围内有 243 个不同的 `addOperations<…>` 符号，demangle 后有 205 个唯一签名（38 个折叠到相同 demangled signature），每个注册 dialect 一个（有些拆成多批）；TPU 路径 registrar 是 `tpu` `0x14aa2c40`（86）、`llo` `0x13e5f940`（325）、`sparse_core` ScDialect `0x14594f60`（115）和 `mosaic_sc` `0x132f9ec0`（1，`RelayoutOp`）。
### 第 2 层 — 每个 slot 委托到 ODS static
每个 `Model<Op>::slotN` thunk 做两件事：定位该 op 的 `Properties` 存储（内联还是离线，由 op flag word 决定），然后 tail-jump 到 ODS 生成的 static。inline-property-offset 惯用法在各 op 中相同：
```asm
; Model<tpu::IotaOp>::getInherentAttr (slot 10 @ 0x14ac18a0)
mov 0x2c(%rdi),%ecx ; load the Operation flag word
shr $0x13,%ecx ; isolate the inline-properties bit
and $0x10,%ecx ; → 0x10 if inline, 0 if out-of-line
lea 0x40(%rdi,%rcx,1),%rsi ; %rsi = &Properties
jmp 0x14b220e0 ; TAIL-JMP to the ODS static, below跳转目标是具体 op 的 ODS static，可恢复为 demangled 符号：
mlir::tpu::IotaOp::getInherentAttr(
mlir::MLIRContext*,
mlir::tpu::detail::IotaOpGenericAdaptorBase::Properties const&,
llvm::StringRef) @ 0x14b220e0
```text
四个 inherent-attribute slot（10–13）都使用这种直接 tail-jmp 形状。
### 第 3 层 — 通过 `UniqueFunction` callback-holder 实现可选 hook
四个*可选* hook，即 `foldHook`（2）、`getCanonicalizationPatterns`（3）、`printAssembly`（7）和 `verifyInvariants`（8），会再多包一层间接调用，形式是 `llvm::UniqueFunctionBase<…>::CallbacksHolder<Op<Op,Traits…>::get<Hook>Fn()::lambda>`。该 slot 加载 holder 指针并执行 `call *(%rax)`（print/parse holder 则是 `call *0x10(%rax)`）。对于 `IotaOp`：
```text
Model<tpu::IotaOp>::foldHook (slot 2 @ 0x14ac1580) → holder @ 0x223413f0
Model<tpu::IotaOp>::printAssembly (slot 7 @ 0x14ac16e0) → holder @ 0x22341400
Model<tpu::IotaOp>::verifyInvariants (slot 8 @ 0x14ac1760) → Op<…>::verifyInvariants @ 0x14ac1de0这一第三层编码了该 op 的完整 trait 和 interface 列表：holder 符号携带整个 Op<OpName, Trait1, Trait2, …, Interface::Trait, …> 模板签名。从 getFoldHookFn holder 符号读取这些模板实参，就是恢复逐 op interface fan-out（下一节）的方法；并且它会与一个独立计数（第二层 interface Model）交叉验证。
陷阱 — 是两层 Model，不是一层。 23-slot
RegisteredOperationName::Model<Op>是 op-registration concept。OpInterfacedispatch（例如“这个 op 有什么副作用？”）通过一个独立的每(op, interface)concept table：mlir::detail::<Iface>InterfaceTraits::Model<Op>，例如MemoryEffectOpInterfaceTraits::Model<mlir::tpu::IotaOp>::getEffects。23-slot Model 的 slot-4hasTrait只是回答“这个 op 是否实现 interface X？”的门禁；interface body 位于第二层 Model 中。重新实现 op dispatch 时必须同时构建二者：这里的 registration vtable，以及逐 interface concept table（后者已为 TPU 路径编目，但其 slot 布局超出本页范围，是第二层 Model 工作的主题）。
调用方汇合点
例如 mlir::Operation::fold @ 0x1d8cd480 会消费该 dispatch：它加载 Impl/Model，然后加载 vptr，再调用 slot 2：
mov 0x30(%rdi),%rdi ; %rdi = OperationName::Impl (= the Model)
mov (%rdi),%rax ; %rax = vptr (address point)
call *0x10(%rax) ; +0x10 from address point = slot 2 = foldHook (@ 0x1d8cd4ad)
```text
---
## tpu Op-Family 分类
86 个 `tpu` op（Model vtable 位于 `0x219aa468 … 0x219aed**` 范围内）会自然地按其 `OpInterface` 签名分组，即每个 op 实现的一组 interface；这些 interface 从该 op 的 `getFoldHookFn` holder 中的 `Op<OpName, Traits…>` 模板实参读取。Interface 缩写为：**MemEff** = `MemoryEffectOpInterface`，**CondSpec** = `ConditionallySpeculatable`，**Bytecode** = `BytecodeOpInterface`，**InferType** = `InferTypeOpInterface`，**OpAsm** = `OpAsmOpInterface`，**RBTerm** = `RegionBranchTerminatorOpInterface`。
```text
#ops OpInterface signature representative ops
---- --------------------------------------------- ----------------------------------------
21 Bytecode, CondSpec, MemEff AllReduce, BroadcastInSublanes, Gather,
Iota, Matmul, MemRefSlice, Pack*, FPToSI,
SIToFP, Transpose, TruncF, Unpack*, …
16 CondSpec, MemEff AssumeLayout, Bitcast, BitcastVreg,
CreateMask, EraseLayout, ExtF, MaskCast,
MemRefBitcast/Reshape/Squeeze, Reshape,
RollVectors, UnrollVectors, Weird, …
13 Bytecode (only) EnqueueDMA, EnqueueIndirectDMA, Log,
LogBuffer, MatmulAccLhs/Pop/PushRhs,
Scan, SemaphoreSignal, Trace*, WaitDMA2
11 (none — pure structural / side-effecting) AllocaSemaphore, Barrier, Delay,
GetBarrierSemaphore, GetInternalScratch,
PRNGRandomBits, PRNGSeed32, Region,
SemaphoreWait, TraceStop, WaitIndirectDMA
10 Bytecode, MemEff Load, ShuffledLoad/Store, Store,
StridedLoad/Store, VectorLoad/Store(Idx)
6 Bytecode, CondSpec, InferType, MemEff AssumeMultiple, Concatenate, DynamicGather,
DynamicRotate, Reciprocal, Rotate
2 CondSpec, InferType, MemEff DeviceId, Relayout
2 Bytecode, InferType GetIterationBound, SublaneShuffle
1 InferType, MemEff FetchAndAddSync
1 CondSpec, InferType, MemEff, OpAsm ScanCount
1 InferType (only) SemaphoreRead
1 Bytecode, CondSpec, MemEff, OpAsm Sort
1 CondSpec, MemEff, RBTerm Yield (the tpu.region terminator)汇总 86 个 op 后，interface fan-out 是：MemoryEffect 59、Bytecode 53、ConditionallySpeculatable 48、InferType 13、OpAsm 2（ScanCount、Sort）、RegionBranchTerminator 1（Yield）。MemoryEffect 计数经过交叉验证：统计独立的第二层 Model MemoryEffectOpInterfaceTraits::Model<mlir::tpu::*> 也得到 59。
这些 op family 可以理解为 TPU 的机器指令类别：MXU matmul（Matmul、MatmulAccLhs/Pop/PushRhs）、DMA（EnqueueDMA、EnqueueIndirectDMA、WaitDMA2、WaitIndirectDMA）、semaphores（AllocaSemaphore、SemaphoreSignal/Wait/Read、GetBarrierSemaphore、FetchAndAddSync）、vector relayout / pack-unpack（Relayout、RollVectors、UnrollVectors、Pack*、Unpack*、BitcastVreg）、memref view ops（MemRefSlice/Bitcast/Reshape/Squeeze、EraseLayout、AssumeLayout）、PRNG（PRNGSeed32、PRNGRandomBits）、conversions（ExtF、TruncF、FPToSI、FPToUI、SIToFP、UIToFP），以及 trace/region structure（Trace*、Region、Yield、Barrier、Delay）。
携带 Properties 的 op
读取全部 86 个 op 的 slot 14（getOpPropertyByteSize），可见 53 个 tpu op 携带非平凡 Properties 结构体（slot 14 返回非零 sizeof；其余 33 个通过 xor %eax,%eax;ret 返回 0）。Properties 大小范围从 8 到 40 字节，多数是 op inherent attribute 的内联存储（例如 IotaOp 的 8 字节 dimensions DenseArrayAttr 句柄，已由 demangled 的 mlir::tpu::IotaOp::getDimensions、readProperties 和 writeProperties static 确认）。
一个特别的9 个 op 子集还额外携带 AttrSizedOperandSegments 可变操作数 segment size 数组（通过 demangled Op<…> 签名上的 AttrSizedOperandSegments trait 验证）：EnqueueDMA、MemRefSlice、SemaphoreSignal、Store、VectorLoad、VectorLoadIdx、VectorStore、VectorStoreIdx、WaitDMA2。最大的 Properties 结构体是 40 字节的 EnqueueDMA 和 Store，随后是 32 字节的 Matmul/Rotate/UnpackSubelements/VectorStore/WaitDMA2；它们把若干 inherent attribute（以及对携带 segment 的 op 而言，大小数组）组合在一个内联 blob 中。
In-dialect rewrite 表面只有 15 个方法
Fold 和 canonicalize slot（2 和 3）在全部 86 个 op 上都存在，但每个都包装一个 no-op，除非该 op 定义了相应方法。按 Op::fold / Op::getCanonicalizationPatterns 的 demangled 符号存在性统计，tpu dialect 的完整 in-dialect rewrite 表面是：
- 6 个真实
fold()：BitcastVregOp、EraseLayoutOp、ExtFOp、MemRefSliceOp、ReshapeOp、TruncFOp。（六个均由符号确认；TruncFOp::fold显示两个符号，即 fold 方法本身，以及执行逐元素转换的内部operator()(APFloat const&, bool&)lambda。） - 9 个真实
getCanonicalizationPatterns()：FPToSIOp、MatmulOp、MemRefBitcastOp、MemRefSliceOp、MemRefSqueezeOp、ShuffledLoadOp、ShuffledStoreOp、UnpackSubelementsOp、UnrollVectorsOp。
tpu dialect 中的其他一切都由 dialect-conversion lowering（tpu → LLO Lowering）进行结构性转换，而不是原地 rewrite。Fold/canonicalize 函数体（例如 MatmulOp::getCanonicalizationPatterns 内部的 rewrite 代数）本页不做反编译，只建立统计（哪些 op 会 rewrite）。
说明 — “86” 修正了朴素
*Op字符串扫描约 157 的过计数。 对mlir::tpu::*Op符号做表面 RTTI 扫描会返回约 157 个命中，但从addOperations<>registrar arity 得到的精确数字是 86 个已注册 op，每个都有 23-slot Model。差值是带*Op后缀但不是已注册 op 的符号，即 Pass 类和 helper（例如LinalgVectorizationPass、MosaicSerdePass、PreCanonicalization、applyLayout），它们会命中朴素扫描，但永远不会到达addOperations，因此没有 Model。编译器概览 出于这个原因刻意不声明tpuop 计数。对于 dispatch-ABI 目的，86 是权威计数。
6,050-Model 统计中的 tpu Dialect
tpu dialect 的 86 个 op 位于全二进制 6,050 个 Model<Op> 实例的统计之中，横跨 60 多个 dialect 命名空间（精确 dialect 计数取决于如何划分 xla 子命名空间，即 xla、xla::cpu、xla::xtile），其中每个 dialect 的 Model 计数都等于它的已注册 op 计数。与 TPU 相关的 dialect 如下：
dialect Models role on the TPU path
------------------ ------ --------------------------------------------------
sparse_core 1471 SparseCore: 115 ScDialect ops + 1356 LlvmTpu tpu_* intrinsics
llo 325 LLO — TensorCore low-level machine ops (the layer below tpu)
LLVM 294 LLVM dialect — SparseCore tail / off the TensorCore path
linalg 99 Mosaic hlo-conversion stage
tpu 86 TPU — TensorCore high-level (Mosaic) target dialect ← THIS PAGE
arith / math / memref 52 / 46 / 32 scalar/math/memref helpers used by tpu, llo, sc
xtile 6 XLA CPU/GPU tile IR (off the TPU device path)
mosaic_sc 1 SparseCore Mosaic — RelayoutOp only
```text
其余 50 多个命名空间（TF 834、spirv 376、ROCDL 350、NVVM 197、HLO 输入语言 mhlo/vhlo/stablehlo/chlo 等）要么是早期就被 legalize 掉的输入语言，要么是路径外的构建残留。对于重新实现者，有两个 TPU 相邻点值得标出：
- **`llo`（325 个 op）** 是直接位于 `tpu` *之下*的 dialect，同一个 23-slot Model ABI 适用。它的大部分（164 个 op）携带 `CondSpec, InferType, MemEff`（scalar/vector ALU），并且只有 `llo::ConstantOp` 定义了真实 `fold()`（LLO constant folder）；没有任何 `llo` op 定义 canonicalizer，因为 LLO 是最底层 MLIR 层，即被 lowering 的层，而不是被 rewrite 的层。
- **`sparse_core`（1471 个 op）** 分解为 115 个高层 `ScDialect` op，加上 **1356 个 `LlvmTpuDialect` `tpu_*` LLVM-intrinsic op**，这是该二进制 MLIR 层中最大的单一 op family；每个都携带 LLVM-op interface（`LLVM::AccessGroup`、`LLVM::AliasAnalysis`），因为每个都会 lower 到一个 LLVM op。它们在 [LlvmTpu Intrinsic Catalog](llvmtpu-intrinsic-catalog.md) 中编目，并由 [LowerToSparseCoreLlvm](lower-to-sparsecore-llvm.md) 消费。它们共享本页记录的完全相同的 23-slot Model ABI。
---
## 置信度汇总
| 断言 | 证据 |
|---|---|
| `Model<Op>` 是 23-slot vtable；全二进制 6,050 个实例 | `_ZTVN4mlir23RegisteredOperationName5ModelI…EEE` 的 `nm` 计数 = 6,050；参考 vtable `Model<xla::PureCallOp>` @ `0x219d4e38` |
| 一个 `tpu` op 上存在全部 21 个 dispatch-slot 符号 | 每个 `Model<mlir::tpu::IotaOp>::<method>`（slot 2–22）都解析到 demangled 符号 |
| Slot 0 是所有 Model 共享的同一个 base dtor | addend `0xfea8820` = `OperationName::Impl::~Impl`，已在 30-Model 分层样本（tpu/llo/sc/mosaic_sc/xla）上确认 |
| `tpu` dialect 恰好有 86 个已注册 op | `Dialect::addOperations<mlir::tpu::AllReduceOp, …>` @ `0x14aa2c40` 有 86 个类型实参 |
| 注册 sink 是 `RegisteredOperationName::insert` | demangled `insert(unique_ptr<OperationName::Impl,…>, ArrayRef<StringRef>)` @ `0x1d8c57a0` |
| Inherent-attr slot tail-jmp 到 ODS static | `Model<IotaOp>::getInherentAttr`（slot 10）tail-jmp → `mlir::tpu::IotaOp::getInherentAttr(MLIRContext*, …IotaOpGenericAdaptorBase::Properties const&, StringRef)` @ `0x14b220e0` |
| 可选 hook 包装在 `UniqueFunction` callback-holder 中 | `Model<IotaOp>::foldHook`/`printAssembly` 加载 holder ptr + `call *(%rax)`/`*0x10(%rax)`；holder @ `0x223413f0`/`0x22341400` |
| `getOpPropertyByteSize` 是内联的 `sizeof(Properties)` | `Model<IotaOp>` = `mov $0x8;ret`；`Model<BarrierOp>` @ `0x14ab1960` = `xor %eax,%eax;ret` |
| `tpu` MemoryEffect fan-out = 59 个 op | `Op<>` trait 扫描 = 59；独立 `MemoryEffectOpInterfaceTraits::Model<mlir::tpu::*>` distinct count = 59 |
| `tpu` in-dialect rewrite 表面 = 6 个 fold + 9 个 canon | 存在 demangled `mlir::tpu::*Op::fold`（6）和 `*Op::getCanonicalizationPatterns`（9） |
| 53 个 `tpu` op 携带非平凡 `Properties`；其中 9 个携带 `AttrSizedOperandSegments` | slot 14 `getOpPropertyByteSize` 在 53/86 个 op 上非零（直接读取：`mov $0xN;ret` 对比 `xor;ret`）；9 个携带 segment 的 op 具有 `AttrSizedOperandSegments` trait；`IotaOp` = 8 字节 `dimensions` property |
| `tpu` op interface 签名（family 分组） | 从每个 op 的 `getFoldHookFn` holder 符号中的 `Op<OpName,Traits…>` 模板实参读取 |
| 朴素 `*Op` 扫描约 157，对比 86 个已注册 op | 86 = registrar arity；约 71 个额外项是未注册 Pass/helper `*Op` 符号，没有 Model |
---
## 交叉引用
- [编译器概览](overview.md) — `tpu` dialect 在五阶段下降和六层 IR 栈中的位置；优化器和 Mosaic 路径的汇合点。
- [MLIR Op-Model 契约](mlir-op-model-contract.md) — 全二进制的 `Model<Op>` ABI 和完整 6,050-Model 统计；该页负责契约，本页是它在 `tpu` dialect 上的具体实例（关联，不重复）。
- [tpu → LLO Lowering](tpu-to-llo-ods.md) — 86 个 `tpu` op 如何被结构性 lowering 到下一层的 325-op `llo` dialect。
- [MHLO → XTile → tpu](mhlo-xtile-tpu-lowering.md) — 产出本页所注册 `tpu` op 的中层 lowering。
- [DialectConversion Legalizer](dialect-conversion-legalizer.md) — lowering 期间消费 `tpu` op 的 conversion-pattern 机制。
- [tpu Program Serialization](tpu-program-serialization.md) — property 存储（slot 18/19，`setPropertiesFromAttr`/`getPropertiesAsAttr`）如何通过 bytecode 往返。
- [LlvmTpu Intrinsic Catalog](llvmtpu-intrinsic-catalog.md) — 共享此 Model ABI 的 1356 个 `sparse_core::tpu_*` intrinsic op。
- [LowerToSparseCoreLlvm](lower-to-sparsecore-llvm.md) — 在 SparseCore 路径上消费 `tpu_*` intrinsic Model 的 pass。
- **二进制：** `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)