RTTI / Vtable 普查

本页所有地址、计数和符号名均适用于来自 libtpu-0.0.40-cp314 wheel 的 libtpu.so：一个 781,691,048 字节的 ELF64 共享对象，build-id 为 89edbbe81c5b328a958fe628a9f2207d（wheel/METADATA/__init__ 版本为 0.0.40；应固定到无歧义的 build-id）。其他 wheel 的每个地址都会不同。

摘要

libtpu.so 发布时未剥离并带完整 RTTI。单是这一事实，就把该二进制中 745 MiB 的静态链接 C++ 变成了一张自描述对象图：每个多态类都留下一个 Itanium-ABI type_info 记录（_ZTI）、一个类型名字符串（_ZTS），并且如果它是具体类，还会留下一个 vtable group（_ZTV）。RTTI 边车保存 160,351 条这样的记录。遍历它们可以在完全不反汇编函数的情况下重建编译器、运行时以及每个 vendored 依赖的整棵 C++ 类森林：继承边位于 type_info 结构中，而 typeinfo↔vtable 绑定把每个类连接到它的分发表。

本页就是这份普查。它确立总计数（160,351 条记录如何拆分为 _ZTI / _ZTV / _ZTS，以及 typeinfo 如何分布到三种 Itanium type_info 风格），按宽度和深度排序主要多态层级（protobuf message 宇宙、两个 MLIR “Operation” 结构、llvm::Pass、xla::HloInstruction、TPU 按 lane-cluster 的树），并记录交叉验证方法：把每个 _ZTI 映射到它的 _ZTV，并从 __class_type_info / __si_class_type_info / __vmi_class_type_info 读出基类链。这个绑定正是 dispatch-table taxonomy 可信的原因：恢复出的分发表只有在附着的类身份正确时才有价值，而这里每个类身份都可对照类型森林检查。

贯穿本页的框架是：这就是 Clang/libstdc++ 发出的 Itanium C++ ABI。熟悉 g++ -frtti 如何布局 vtable group 的读者，也就是 [offset-to-top][typeinfo-ptr][fn0][fn1]…] 且 address point 位于 +0x10，已经知道数据模型。唯一的转折是，该二进制被预链接成类似非 PIE 的镜像：文件字节中的 typeinfo 指针、基类指针和函数指针均为零，由 R_X86_64_RELATIVE 重定位 addend 提供，因此继承图是从重定位表而不是原始 .data.rel.ro 镜像中重建的。

对于重新实现，也就是从二进制重建这份普查，契约是：

记录分类法： 160,351 条 RTTI 记录如何分区为 typeinfo 结构、vtable group 和名称字符串，以及 typeinfo 的第一个 word 如何识别其 type_info 风格。
边恢复规则： 在基类指针是重定位 addend 的前提下，如何从 __si_class_type_info+0x10 读出单个基类，并从 __vmi_class_type_info+0x18 读出基类数组。
typeinfo↔vtable 绑定： _ZTV+8 → _ZTI 不变量，多继承 secondary sub-vtable 如何重复本类 typeinfo，以及 “abstract = 有 _ZTI 但没有 _ZTV” 在结构上意味着什么。


RTTI 边车记录	160,351 (`_ZTI` 60,457 · `_ZTV` 39,244 · `_ZTS` 60,650 · 2 demangler-prefix strings)
`type_info` 风格	`__class_type_info` 16,754 · `__si_class_type_info` 42,447 · `__vmi_class_type_info` 680
继承边	~43,807 (42,440 single-base + 1,367 multi/virtual-base)
类根 / 真实层级	16,761 roots · 2,094 with ≥2 descendants
最宽层级	`proto2::MessageLite` — 8,013 descendants, depth 3 (`_ZTI` `0x22034138`)
最深层级	`grpc::Service` — depth 10 (`_ZTI` `0x216162d8`, `_ZTV` `0x216162e8`)
Vtable↔typeinfo 绑定	39,243 / 39,244 vtables bind to own-class `_ZTI`; 0 cross-class; 1 RTTI-less (`libunwind`)
重定位来源	`.rela.dyn` @ file `0x9170`, 1,069,006 `R_X86_64_RELATIVE` addends

普查概要

记录拆分

RTTI 边车是一个包含 160,351 条记录的数组，每条记录都是从 .rodata / .data.rel.ro 收集的 (struct-address, mangled-symbol, name-string) 三元组。按 mangled symbol 的三字母 Itanium 前缀分组，可得到类型系统的顶层形状：

记录种类	前缀	数量	含义
Typeinfo struct	`_ZTI`	60,457	每个多态类型一个（class/si/vmi + pointer/pbase/fundamental）
Vtable group	`_ZTV`	39,244	每个具体多态类一个（至少有一个可实例化对象）
Typeinfo-name string	`_ZTS`	60,650	`_ZTI` 指向的可读 mangled 类型名
Demangler-prefix strings	(none)	2	二进制内 demangler 使用的字面量 `"typeinfo for "` / `"typeinfo name for "`

算术完全闭合：60,457 + 39,244 + 60,650 + 2 = 160,351。这三个计数是彼此独立的测量：_ZTI 和 _ZTV 是 .data.rel.ro 中的结构地址，_ZTS 是 .rodata 中的字符串地址。它们内部一致：typeinfo 多于 vtable（60,457 对 39,244），正是因为 59,881 个 class-kind typeinfo 中有 20,638 个是抽象的，也就是带身份但不拥有 vtable 的 pure-virtual interface base（见 § 交叉验证）。_ZTS 字符串略多于 _ZTI 结构（60,650 对 60,457），因为少数名称字符串被共享，或为 _ZTI 被折叠的类型发出。

NOTE — 160,351 这个数字是记录计数，不是不同类计数。一个 C++ 模板实例化 N 次会产生 N 对不同的 _ZTI/_ZTS（例如每个 std::__shared_count<...> 特化、每个 RegisteredOperationName::Model<Op>），因此群体被模板爆炸主导。“不同逻辑类”数量要小得多；本普查按 ABI 发出的实例化类型计数。
Note (count provenance): 权威 RTTI 计数来自去重符号表，而不是结构遍历或反编译树 grep。因为 libtpu.so 发布时未剥离，每个 _ZTI/_ZTV/_ZTS 都带一个本地 d 符号，所以 nm 是穷尽的，任何结构遍历都不可能找到比符号更多的记录。直接对二进制运行 nm 得到 60,457 个 _ZTI、39,244 个 _ZTV、60,650 个 _ZTS（总计 160,351）。更高的数字，例如 160,566（60,471 / 39,246 / 60,847），是过计数，这里不使用。（复现：nm libtpu.so | rg -c '_ZTI'，等等。）

按命名空间的群体

按每个 typeinfo 的 mangled _ZTIN<len><namespace> 符号中的前导命名空间 token 分桶，可得到类型森林的人口统计，并确认主导多态表面的，是编译器/运行时核心，而不是 vendored support libraries。（这些是前导命名空间桶，不是 demangled 顶层 token 计数：后者会把每个 absl::StatusOr<xla::…> / std::unique_ptr<…> wrapper 都过度归功于 xla/tensorflow，即使其 owning class 位于别处。完整分解由 namespace census appendix 负责。）

命名空间	Typeinfos	角色
`mlir`	13,091	MLIR dialect、op、pass、pattern、interface
`asic_sw`	11,379	TPU driver / profiler / ISA event-control 树
`tensorflow`	3,108	TF op-kernel、graph、TFRT TPU
`xla`	3,036	HLO、thunk、HLO pass、jellyfish codegen
`llvm`	2,940	嵌入式 LLVM 后端（codegen、Attributor、VPlan）
(anon ns)	2,352	翻译单元局部类（`_GLOBAL__N_`）
`dnnl`	1,888	oneDNN primitive descriptor / JIT kernel
`std`	1,787	libstdc++ 模板管道
`grpc_core`	1,502	gRPC ref-counted / orphanable object 树
`platforms_deepsea`	576	TPU ISA encoder、debugger service

在 60,457 个 _ZTI 结构中，46,078 个携带前导命名空间，14,379 个是全局作用域或复合类型（pointer/function/substitution）。mlir + xla + llvm + tensorflow + asic_sw + platforms_deepsea 的总和是编译器/运行时核心；std、absl、dnnl、Eigen、grpc、OR-tools 是 vendored support。不过，最大的单棵树位于 proto2 下（protobuf message 宇宙）：宽而浅，并不是最大的命名空间。

顶层层级

一个层级的宽度是其传递后代数量；深度是从根向下的最长链。森林有 16,761 个类根，其中 2,094 个是真实层级（≥ 2 个后代）。下表是整理后的顶层梯队，即结构上居中的树；纯模板/管道树（例如宽度 1,338 的 std::__shared_count，四个宽度各 376 的 mlir::spirv::Query*InterfaceTraits::Concept）被排除，除非它们就是该二进制的定义性结构。

宽度	深度	根类	Root `_ZTI`	Root `_ZTV`	角色
8013	3	`proto2::MessageLite`	`0x22034138`	abstract	protobuf message 反射（Message + 8,007 个生成消息）
6142	9	`mlir::Pattern`	`0x21cea698`	abstract	MLIR rewrite / conversion / lowering pattern 树
6052	2	`mlir::OperationName::InterfaceConcept`	`0x217b1000`	abstract	MLIR op-interface `Model<Op>` dispatch
2069	6	`dnnl::impl::c_compatible`	`0x21b69258`	abstract	oneDNN primitive-descriptor / JIT 基类
1122	4	`tensorflow::OpKernel`	`0x218114c8`	`0x218113d8`	TF op-kernel 基类（TPU embedding / XLA kernel）
821	1	`asic_sw::…::profiler::EventControlInterface`	`0x2175c798`	abstract	按 lane-cluster 的 TPU perf-counter event control
628	5	`llvm::Pass`	`0x21ced3b8`	`0x21ced328`	LLVM legacy pass-manager 层级
606	7	`mlir::Pass`	`0x21c2c450`	`0x21c2c3d8`	MLIR pass 层级（Op/Module/Function pass）
551	5	`Xbyak::CodeArray`	`0x21b6d738`	`0x21b6d818`	Xbyak JIT code-generator 基类（oneDNN 后端）
442	6	`grpc_core::PolymorphicRefCount`	`0x21ca0128`	abstract	gRPC ref-counted object 基类
361	4	`xla::HloPassInterface`	`0x217f4428`	abstract	XLA HLO compiler-pass interface
329	9	`llvm::AbstractState`	`0x218540e8`	abstract	LLVM Attributor abstract-attribute 状态机
299	8	`llvm::IRPosition`	`0x21b345d8`	abstract	LLVM Attributor IR-position 抽象
299	8	`llvm::AADepGraphNode`	`0x21854120`	`0x218540f8`	LLVM Attributor dependency-graph 节点
184	4	`grpc_core::Orphanable`	`0x21ca01b8`	abstract	gRPC orphanable object 基类
166	3	`llvm::cl::Option`	`0x21fe0cc0`	`0x21fe0c58`	LLVM command-line option 基类（供给 MLIR `PassOptions`）
140	4	`tsl::core::RefCounted`	`0x215f9b18`	abstract	TSL/TF ref-counted 基类
114	6	`riegeli::Object`	`0x220291a8`	`0x22029168`	riegeli record-IO object 基类
100	3	`mlir::DialectInterface`	`0x21cea480`	abstract	MLIR dialect-interface 基类
68	4	`xla::HloInstruction`	`0x21d2ce88`	`0x21d2cde8`	XLA HLO instruction 树
53	3	`xla::cpu::Thunk`	`0x219b15a0`	`0x219b1550`	XLA:CPU runtime thunk 树
44	10	`grpc::Service`	`0x216162d8`	`0x216162e8`	gRPC generated-service 基类（二进制中最深）
23	3	`llvm::MCStreamer`	`0x21b60b70`	`0x21b60630`	LLVM MC streamer 树（+ TPU SparseCore `SCStreamer`）

GOTCHA — 最宽的若干棵树中有四棵是抽象的：它们的 _ZTV 列显示 “abstract”，因为根带有 _ZTI 但不拥有 vtable。proto2::MessageLite、mlir::Pattern、mlir::OperationName::InterfaceConcept、HloPassInterface、AbstractState、IRPosition、EventControlInterface 都是 pure-virtual interface base。寻找 “the MessageLite vtable” 来 hook 的重新实现者找不到它：dispatch 位于下一层，存在于每个具体子类自己的 _ZTV 中。这是反复出现的按 generation family 模式：接口根加上一片平坦的具体叶子，每个叶子都完全 override（见 TPU codec / cycle-table / encoder family，按代函数调度器）。

这份排名有四个结构事实值得强调，因为读者应该从中学到它们，而不是死记：

protobuf message 宇宙是最大的单棵树。 proto2::MessageLite → proto2::Message（8,007 个后代）→ 每个 .proto 类型一个具体生成消息。它宽（8,013）但浅（深度 3）：protobuf 生成类是一片平坦扇出，而不是深链。这是该二进制嵌入的 FileDescriptorProto 记录的内存中反射支撑。

MLIR 有两个不同的 “Operation” 结构，而且两者都不是多态的 mlir::Operation。 mlir::Operation 本身是非多态的：它没有 _ZTI，也没有 vtable；MLIR op 完全不携带 C++ virtual table。Op 行为通过两种方式调度，两者都表现为巨大的 RTTI 树：

text

(a) Type-erased op-interface dispatch  (the "op model"):
    mlir::OperationName::InterfaceConcept   (6,052 desc, depth 2)
      └─ mlir::RegisteredOperationName
           └─ Model<Op>   (one per registered op — verify/parse/print/fold)

(b) The rewrite / lowering pattern tree:
    mlir::Pattern   (6,142 desc, depth 9 — second-widest tree)
      └─ RewritePattern
           └─ ConversionPattern
                └─ ConvertToLLVMPattern
                     └─ ConvertOpToLLVMPattern
                          └─ ... (TPU SparseCore lowering is the deepest branch)

Model<Op> 数组正是 dispatch-table taxonomy 中的 MLIR-Op-Model 类；Pattern 树是整个二进制中第二宽的树。

LLVM 后端被完整嵌入，并且是最深的 codegen 结构。 llvm::Pass（628 个后代）拆分为经典 legacy-PM 类别；主分支是 FunctionPass（506）→ MachineFunctionPass（351，其中 324 个 direct）→ AMDGPU / PPC / TPU machine passes。除 TPU target 外，libtpu.so 还嵌入完整的 AMDGPU、PPC、ARM、AArch64 和 X86 后端，每个 *TargetLowering 都存在于 TargetLoweringBase 下。new-PM Attributor 贡献两条并行的 8–9 层深 CRTP 模板链（AbstractState 深度 9，IRPosition/AADepGraphNode 深度 8），是该二进制中仅次于 gRPC service tree 的最深 CRTP 链。

TPU 特有树宽而平，按 lane cluster 分区。 EventControlInterface（821 个后代，深度 1，全部是 direct leaves）是按 lane-cluster 的性能计数器 event-control 层级，其按 cluster 拆分方式与 S-ISA event table 相同。四棵 TPU HAL 树（TpuCodec、isa::Encoder、CycleTable，以及运行时 TpuHal/TpuChip/TpuCore 树）都是 pure-virtual interface root，其按代叶子各自完全 override，这是按代 vtable-family dispatch 而不是 switch 的结构签名。

`xla::HloInstruction` — 一个完整树示例

HLO instruction 树是规范的中型层级，也是最可能被 hook 的层级，因此值得给出完整结构。根 _ZTI 0x21d2ce88，根 _ZTV 0x21d2cde8；38 个 direct subclass 上共有 68 个传递后代；9 个 internal（多层）节点，59 个具体叶子；深度 4。

内部节点	children	flavor	vtable
`HloDimensionsInstruction`	7	si	`0x21d2fb20`
`HloCollectiveInstruction`	5	si	`0x21d2dd40`
`HloSendRecvInstruction`	4	si	`0x21d2d9f8`
`HloChannelInstruction`	3	si	`0x21d2d8b0`
`HloCallableInstruction`	3	vmi	`0x21d2ea00`
`HloBatchNormInstruction`	3	si	`0x21d2d1b0`
`HloDynamicIndexInstruction`	2	si	abstract
`HloAllReduceInstructionBase`	2	si	`0x21d2de90`
`HloAsyncInstruction`	1	si	`0x21d2d4d0`

最深链（深度 4）是 HloInstruction → HloChannelInstruction → HloCollectiveInstruction → HloAllReduceInstructionBase → HloAllReduceInstruction。成对的 xla::DfsHloVisitorBase<HloInstruction const*>（52 个后代，_ZTI 0x21d2c790）是 double-dispatch visitor 对应物。HloCallableInstruction 是树中唯一的多继承节点，是一个 __vmi 类，混合 HloInstruction（0x21d2ce88）和空的 HloAliasible mixin（0x21d2ce98）；因为第二个基类为空，它不会产生 secondary sub-vtable，所以 0x21d2ea00 处的 __vmi group 仍然只在 +8 携带一个本类 typeinfo 指针。

RTTI 记录种类

每个 _ZTI 都是 Itanium ABI 的某种 type_info 风格，其风格编码在 typeinfo 结构的第一个 word 中：它是指向元类型 vtable（type_info 子类自身的 vtable）的 address point 的指针，而元类型 vtable 只来自一个很小的固定集合。因此，对 _ZTI+0 处的 word 做直方图即可在不解析名称的情况下分类全部 60,457 个 typeinfo：

Metatype vtable @`_ZTI+0`	数量	`type_info` 风格	基类信息
`0x220485b0`	16,754	`__class_type_info`	none — a root with no base
`0x22048600`	42,447	`__si_class_type_info`	single public base @ `+0x10`
`0x22048668`	680	`__vmi_class_type_info`	base array @ `+0x18`, count @ `+0x14`
`0x220486d0`	282	`__pointer_type_info`	— (not a class hierarchy)
`0x22048528`	261	`__function_type_info`	— (not a class hierarchy)
`0x22048560`	22	`__enum_type_info`	—
(tail)	11	`__fundamental_type_info` (10, `0x220483d8`) + `__pointer_to_member_type_info` (1, `0x22048708`)	—

Class-kind typeinfo（class + si + vmi）总计 59,881；576 个 pointer/function/enum/fundamental 记录不是类层级，因此正确地作为森林根保留。

QUIRK — _ZTI+0 处的直方图值是元类型 vtable 的 address point，它位于 vtable group base 上方 0x10。__class_type_info 符号（'vtable for'__cxxabiv1::__class_type_info）位于 0x220485a0，但写入每个 __class_type_info typeinfo 第一个 word 的值是 0x220485b0 = 0x220485a0 + 0x10。同理，__si 是 0x220485f0 → 0x22048600，__vmi 是 0x22048658 → 0x22048668。把直方图值与符号地址匹配的重新实现者会发现每个分类都 “off by 16”，除非加上 address-point 偏移。这与支配普通 vtable 的 +0x10 address-point 约定相同（offset-to-top 在 +0，typeinfo 在 +8，第一个 method 在 +0x10）。

读取基类链

三种 class 风格的差异仅在于它们如何编码父类。布局如下，在每个 _ZTI 处读取（64-bit）：

text

common header (all type_info):
  +0x00  metatype vtable ptr (+16)   -> identifies the flavor (table above)
  +0x08  ptr to the _ZTS name string

__si_class_type_info  (single inheritance):
  +0x10  ptr to the single base _ZTI         -> exactly ONE edge

__vmi_class_type_info (multiple / virtual inheritance):
  +0x10  u32 flags
  +0x14  u32 base_count
  +0x18  base_count x { ptr base_ZTI ; long offset_flags }   -> base_count edges
            offset_flags: (val >> 8)  = sub-object byte offset
                          bit0        = base is virtual
                          bit1        = base is public

// Edge recovery — re-derives the 43,807-edge inheritance forest.
// Crucial caveat: the base pointers are 0 in the file image; their real
// values are R_X86_64_RELATIVE relocation ADDENDS. Read the addend at the
// vaddr, not the file byte.
function recover_edges(zti_addr, flavor):
    switch flavor:
      case __class:                                  // 16,754 records
          return []                                  // root, no parent

      case __si:                                     // 42,447 records
          base = reloc_addend_at(zti_addr + 0x10)    // ONE base ZTI
          if base is a known _ZTI:
              emit edge  derived=zti_addr -> base    // 42,440 valid (+7 dangling)

      case __vmi:                                    // 680 records
          n = read_u32(zti_addr + 0x14)              // base_count (plain bytes)
          for i in 0 .. n-1:
              base    = reloc_addend_at(zti_addr + 0x18 + 16*i)
              off_flg = read_long  (zti_addr + 0x18 + 16*i + 8)
              emit edge  derived=zti_addr -> base
              record sub_object_offset = off_flg >> 8     // Table D values
              record is_virtual        = off_flg & 1
              record is_public         = off_flg & 2
          // base-count distribution: 1->121  2->503  3->34  4->4  7->18

GOTCHA — _ZTI+0x10/+0x14 处的 flags 和 base_count 是普通的 .data.rel.ro 字节（无重定位），但数组中的基类指针是重定位 addend。把整个 __vmi 记录当文件字节读取的重新实现者会得到正确的 base_count，以及每个基类指针的 0。.data.rel.ro 映射有固定的 vaddr→file-offset delta 0x200000（vaddr 0x215f81a0 ↔ file 0x213f81a0）；实际边来自 .rela.dyn 中的 1,069,006 个 R_X86_64_RELATIVE 条目。

核心中的多继承

680 个 __vmi 类很少见（class typeinfo 的 1.1%），但结构上很重要。它们的 base-count 分布是 1→121, 2→503, 3→34, 4→4, 7→18。121 个单基类 __vmi 记录之所以是 __vmi，只是为了携带非零 sub-object offset；18 个七基类记录是 abseil raw_hash_set policy CRTP 管道，不是编译器逻辑。真正的编译器核心 diamond 及其解码出的 sub-object offset 是：

派生类	bases	sub-object offsets (`offset_flags >> 8`)
`llvm::TPUScheduleDAGSlackModulo`	4	`ScheduleDAGInstrs` +0 · `TPUScheduleDAGModulo` +3424 · `TPUScheduleDAGComposeFifo` +4528 · `ScheduleDAGInstrsWrapper` +4672
`llvm::TPUScheduleDAGSwingModulo`	3	`TPUScheduleDAGSwing` +0 · `TPUScheduleDAGModulo` +3728 · `ScheduleDAGInstrsWrapper` +4832
`llvm::RABasic`	3	`MachineFunctionPass` +0 · `RegAllocBase` +56 · `LiveRangeEdit::Delegate` +784 (non-public)
`llvm::VPRecipeBase`	3	`VPDef` +0 · `ilist_node_with_parent` +16 · `VPUser` +32
`llvm::legacy::FunctionPassManagerImpl`	3	`Pass` +0 · `PMDataManager` +32 · `PMTopLevelManager` +416
`xla::TpuHostTransferManager`	3	`TransferManager` +0 · `PjRtHostTransferManager` +8 · `enable_shared_from_this<…>` +72

NOTE — 每个真正的编译器核心 diamond 都解码为 offset_flags & 1 == 0，也就是核心中全是非虚继承。MLIR/XLA/LLVM/TPU 代码中没有 virtual base，这意味着 construction vtable（_ZTC）和 VTT（_ZTT）在核心中基本不存在；只有 abseil/std 管道尾部可能携带它们。除 RABasic → LiveRangeEdit::Delegate 外，bit1 public flag 处处置位（delegate mixin 是 non-public）。整体上最常见的 __vmi 模式是 mlir::detail::PassOptions::Option = {llvm::cl::opt, PassOptions::OptionBase}，即 MLIR pass 命令行选项适配器，出现数百次。

交叉验证：RTTI ↔ Vtable {#cross-validation-rtti--vtable}

普查的可信度建立在一个把类型系统两半连接起来的 ABI 不变量上：vtable group 会在偏移 +8 存储同一个精确类的 typeinfo 指针。 group 布局为

text

Itanium vtable group (64-bit), read at each _ZTV symbol:
  +0x00  offset-to-top (long; 0 for the primary sub-vtable, signed for secondaries)
  +0x08  typeinfo pointer  -> the class's own _ZTI        <== THE BINDING
  +0x10  address point: virtual-function slot 0, slot 1, ...

For a multiple-inheritance class the group repeats:
  [ primary  sub-vtable: off-to-top=0,  _ZTI(own), fns... ]
  [ secondary sub-vtable per non-primary polymorphic base:
        off-to-top<0, _ZTI(own again), thunk-fns... ]

交叉验证会遍历每个 _ZTV，读取 _ZTV+8 处的重定位 addend，并确认它落在同一类的 _ZTI 上（join key = _ZTV/_ZTI/_ZTS 三元组共享的 mangled class suffix）。结果具有决定性：

数量	计数
Vtable groups (`_ZTV`)	39,244
Vtable → own-class `_ZTI` at `+8` (direct)	39,185
Vtable → own-class `_ZTI` via MI secondary structure	58
Vtable → own-class `_ZTI` (total)	39,243
Vtable → different-class `_ZTI` (mismatch)	0
Vtable with no `_ZTI` anywhere (RTTI-less)	1 (`libunwind`)
Concrete class typeinfos (own a vtable)	39,243
Abstract class typeinfos (no vtable)	20,638

39,244 个 vtable 中有 39,243 个绑定到自己的类 typeinfo，且跨类 mismatch 为零；该绑定经过双向验证。唯一例外有完整解释：

QUIRK — 唯一没有 RTTI 的 vtable 是 _ZTV 0x22048990 处的 libunwind::UnwindCursor<LocalAddressSpace, Registers_x86_64>（一个 16 槽 group）。它的 +8 typeinfo 槽为 NULL，且该类没有 _ZTI 符号，因为 libunwind 是用 -fno-rtti 编译的。这是 no-RTTI 翻译单元的正确行为，不是绑定缺陷，而且它是 745 MiB 二进制中唯一这样的单元。libtpu.so 中其他每个 vtable 都携带干净的 typeinfo 指针。

绑定证明什么

三个结构检查把 +8 不变量转化为层级交叉检查：

单继承单调性。 对于双方都拥有 vtable 的 __si derived/base 对，derived vtable 必须至少与 base vtable 一样大，因为单继承会把新的 virtual slot 追加到 base 前缀之后。在 2,000 对随机样本上，derived_vtable_size ≥ base_vtable_size 2,000/2,000 成立（0 个 violation）。RTTI 基类边和 vtable 布局一致。

多继承 secondary sub-vtable。 在 509 个拥有 vtable 的 __vmi 类中，508 个至少包含一个本类 typeinfo 指针；7 基类 gRPC EventMgrEventEngine 在其 0x1d0 字节 group 内恰好有 7 次本类 typeinfo 重复（每个 sub-object 一次）。唯一的 “0-secondary” 情况是空 mixin 模式（HloCallableInstruction：基类 HloInstruction + 空 HloAliasible → 一个 sub-vtable，一个位于 +8 的本类 typeinfo），这是正确的 ABI 行为。

抽象根确认。 宽度排名中标记为 “abstract” 的全部 16 个抽象接口根：MessageLite（0x22034138）、mlir::Pattern（0x21cea698）、OperationName::InterfaceConcept（0x217b1000）、EventControlInterface（0x2175c798）、PolymorphicRefCount（0x21ca0128）、HloPassInterface（0x217f4428）、AbstractState（0x218540e8）、IRPosition（0x21b345d8）、Orphanable（0x21ca01b8）、tsl::core::RefCounted（0x215f9b18）、DialectInterface（0x21cea480）、OffloadFactory（0x218fffd8）、ErrorInfoBase（0x21fe1a70）、isa::Encoder（0x21cb6a20）、TpuCodec（0x21d35858）、CycleTable（0x21c20008）都被独立确认不拥有同类 vtable，验证了 abstract-vs-concrete 拆分。

普查如何交叉检查 dispatch taxonomy

typeinfo↔vtable 绑定是本页与 Dispatch-Table Taxonomy 之间的桥梁：taxonomy 恢复出的每个 dispatch-table 都是一个 vtable group，而它的 +8 指针现在可以确定地命名其类。每个顶层层级都映射到一个 taxonomy 类，按代 TPU family 则解析为具体叶子 vtable：

TPU family root (abstract)	Concrete leaf vtables
`tpu::TpuCodec` (`0x21d35858`)	Jellyfish `0x21d360f0` · Dragonfish `0x21d35800` · Pufferfish `0x21d36148` · Viperfish `0x21d361a0` · Ghostlite `0x21d35c00`
`xla::jellyfish::CycleTable` (`0x21c20008`)	Jf `0x21c1ffb8` · Pf `0x21c20048` · Vf `0x21c200c8` · Glc `0x21c20148` · Gfc `0x21c201c8`
`platforms_deepsea::…::isa::Encoder` (`0x21cb6a20`)	EncoderJf `0x21d36ca8` · EncoderDf `0x21d36be0` · per-family `EncoderGl`/`EncoderVf`/`EncoderPf*` leaves

每个 codec 叶子都是没有中间层的 __si 单继承类；连续的 CycleTable 叶子地址（0x21c1ffb8…0x21c201c8）是 dispatch taxonomy 枚举为 cycle-table family 的按代 cost-model table 的 group base。模式是统一的：一个没有自己 vtable 的抽象接口根，以及一片平坦的按 codename 具体叶子，每个叶子都完全 override，这正是按代函数调度器中记录的按代 vtable-family dispatch。

层级深度分布

在 2,094 个真实层级中，深度强烈偏向浅树：二进制中的大多数多态是一两层 override，而不是深继承：

text

depth  1: 1,615    depth  6:   4
depth  2:   264    depth  7:  10
depth  3:   136    depth  8:   3
depth  4:    46    depth  9:   4
depth  5:    11    depth 10:   1

唯一的深度 10 层级是 grpc::Service → 七层堆叠的 TpuDebugService::WithAsyncMethod_* gRPC 模板 wrapper → TpuDebugServiceImpl。33 个层级有超过 100 个传递后代。该形状确认二进制的多态由宽而平的接口扇出主导（按代 family、protobuf 和 MLIR-op 扇出），深链局限于模板密集的 CRTP 代码（Attributor、gRPC service wrapper）。

本普查未解析的内容

逐槽 virtual-method 标签。 绑定证明了每个 vtable 的 address point 和 typeinfo 身份，但没有把每个函数指针槽（_ZTV+0x10+8i addend → symbol）解析成 method 名称。该槽位遍历是下一层；其他页面只对四棵 TPU HAL 树执行了它。
模板实例化折叠。 每个 *<...> 实例化都计为一个不同类型；宽度排名会把 std::__shared_count（1,338）和四个 spirv::Query*::Concept（各 376）计为独立树。整理后的顶层梯队有意选择非模板根，以免它扭曲结果。
非多态类。 mlir::Operation、mlir::Value、TpuHalCommonStates 以及类似类没有 typeinfo，对 RTTI 遍历不可见；它们的布局必须来自成员访问分析，而不是 RTTI。
7 个悬空 __si 基类指针（base addend 不是已知 _ZTI）以及 __vmi 2-base 计数中的 2 类漂移尚未分流；两者都小于 0.02%，不影响任何顶层层级结果。

交叉引用

二进制取证概览 — 文件级框架：section table、符号计数，以及本普查遍历的 745 MiB 形状。
调度表分类 — 负责表计数分类；本普查为每个恢复的 dispatch table 提供经过验证的类身份。
按代函数调度器 — 抽象 TPU-family root 解析到的按 codename 叶子 vtable dispatch 模式。
多态调度入口点 — 在调用点消费这些绑定的 indirect-call sites 和 top-vtable classes。

RTTI / Vtable 普查 ​

摘要 ​

普查概要 ​

记录拆分 ​

按命名空间的群体 ​

顶层层级 ​

xla::HloInstruction — 一个完整树示例 ​

RTTI 记录种类 ​

读取基类链 ​

核心中的多继承 ​

交叉验证：RTTI ↔ Vtable {#cross-validation-rtti--vtable} ​

绑定证明什么 ​

普查如何交叉检查 dispatch taxonomy ​

层级深度分布 ​

本普查未解析的内容 ​

交叉引用 ​

RTTI / Vtable 普查

摘要

普查概要

记录拆分

按命名空间的群体

顶层层级

`xla::HloInstruction` — 一个完整树示例

RTTI 记录种类

读取基类链

核心中的多继承

交叉验证：RTTI ↔ Vtable {#cross-validation-rtti--vtable}

绑定证明什么

普查如何交叉检查 dispatch taxonomy

层级深度分布

本普查未解析的内容

交叉引用