Internal Wiki

主题

多态分派入口点 ​

本页中的所有地址均适用于 libtpu-0.0.40-cp314 wheel 中的 libtpu.so（build-id 89edbbe81c5b328a958fe628a9f2207d）。每个 VA 都是未重定位镜像中的加载地址；可执行段满足 VA == file-offset。其他构建会有所不同。

摘要 ​

这是控制流扇出导航页。一个 745 MB 的 stripped C++ 二进制不是由 call sub_X 边组成的树，它的热路径会经过间接调用，其中目标是从 vtable 槽位或类型擦除 callable 中读取出来的函数指针。若要从 PJRT 入口表面一路跟踪执行到 TPU ISA 编码器，逆向工程师必须站在每个这样的“枢纽”处，并把 call qword ptr [reg+0xNN] 解析为具体实现者。本页映射重要的十几个枢纽，把每个枢纽锚定到其调用点地址和它读取的 vtable 槽位，然后说明通过把槽位索引连接回 RTTI 普查来解析任意间接调用的一般流程。

这个二进制中有两种 C++ 分派形态，外加第三种并非 C++ vtable 的形态。(a) vtable 槽位分派 — mov (obj),%vptr ; call *0xNN(%vptr) — 占绝大多数；槽位索引是 0xNN / 8。(b) 函数指针分派 — call *%reg — 被调用者先被加载到寄存器中；这是 mlir::PatternApplicator、llvm::function_ref 以及 std::function/AnyInvocable 池的分派方式，而只 grep call *[reg+off] 的导航者会完全漏掉它。(c) PJRT C-ABI 表面是一个扁平函数指针结构体（PJRT_Api），一个在首次调用时填充一次的 C 分派表；它在结构上不同于 C++ vtable，并且要通过读取结构体初始化器来解析，而不是通过 RTTI。

机制因 IR 层而异，而这种差异正是重点。XLA HLO 层分派经典的访问者模式：HloInstruction::Visit 是一个 132 路 opcode switch，其中每个 case 都尾调用 DfsHloVisitor vtable 的不同 Handle<Opcode> 槽位。MLIR 并不为每个 op 使用一个 vtable，而是使用基于 concept 的 Op Model，即每个 op 生成的分派对象，其 foldHook/hasTrait 槽位由 mlir::Operation::fold 间接读取。pass manager（HloPassInterface::Run、OpToOpPassAdaptor::run）是薄 trampoline，会尾跳到单个槽位。TPU codegen（CodeGenerator::EmitInstruction）在按硬件代际填充的 152 槽 IsaEmitter vtable 中跨 81 个槽位扇出。

用于导航的契约是：

  • 两种半分派形态以及如何在反汇编中识别每一种。
  • 主要枢纽：调用点地址、vtable VA + 它读取的槽位，以及该槽位解析到什么。
  • 解析流程：槽位索引 -> 通过 RTTI vtable 普查得到候选实现者；如何处理 RTTI 不覆盖的 C-ABI 和函数指针形态。
全二进制最热偏移0x10（槽位 2 — 每个层级中的第一个非 dtor virtual：name()/Compute()/foldHook）
HLO 访问者扇出HloInstruction::Visit（0x1e585660）— 进入 DfsHloVisitor vtable 的 132-case opcode switch
访问者 vtable_ZTVN3xla17DfsHloVisitorBaseIPNS_14HloInstructionEEE @ 0x21d2c320（address point +0x10）
HLO pass trampolineHloPassInterface::Run（0x1e472a60，6 B）— jmp *0x28(%rax) = 槽位 5
MLIR pass 主体OpToOpPassAdaptor::run（0x1cb6dc20）— call *0x38(%rax) = 槽位 7（runOnOperation）
MLIR Op-Model foldmlir::Operation::fold（0x1d8cd480）— call *0x10(Model) = 槽位 2（foldHook）
PJRT C-ABI来自 GetTpuPjrtApi（0xe6aa440）的扁平 PJRT_Api 结构体 — C 函数指针，不是 vtable

一览：分派枢纽 ​

每个枢纽都是导航者会反复遇到的一个间接调用点。slot = off / 8。槽位标签是 RTTI 普查中的逐槽位名称；“扇出”是单个站点可到达的不同实现者数量。

枢纽调用点 VA经由（vtable / 形态）槽位 / off扇出
HLO 访问者 opcode 分派0x1e585660（Visit）DfsHloVisitor @ 0x21d2c320132 个不同槽位，off 0x20-0x438每个 opcode 一个 Handle<Op>
HLO 每节点 pre-hook0x1e5866e0（PostOrderDFS）同一 vtable137 / 0x448（Preprocess）按具体 visitor
HLO 每节点 post-hook0x1e5866e0同一 vtable138 / 0x450（Postprocess）按具体 visitor
HLO 每节点 gate0x1e5866e0同一 vtable139 / 0x458（ShouldProcessNode）按具体 visitor
HLO finish-hook0x1e584660（Accept）同一 vtable136 / 0x440（FinishVisit）按具体 visitor
HLO pass 主体0x1e472a60（Run）HloPassInterface5 / 0x28（RunImpl）每个 HLO pass
HLO pass 主体（uptr）0x1e472a80（Run）HloPassInterface6 / 0x30（RunImpl uptr）每个嵌套 pipeline
MLIR pass 主体0x1cb6dc20（OpToOpPassAdaptor::run）mlir::Pass7 / 0x38（runOnOperation）每个 MLIR pass
MLIR Op-Model fold0x1d8cd480（Operation::fold）Op Model concept2 / 0x10（foldHook）每个已注册 MLIR op
CPU thunk execute0x1c0f0320（TracedExecute）xla::cpu::Thunk5 / 0x28（Execute）每种 thunk
TPU ISA emit0x14043a40（EmitInstruction）IsaEmitter（152 槽）81 个槽位，off 0x50-0x490按代际 {Pf,Vf,Gl,Gf} emitter
TpuHal 硬件启动0x1e811ea0（InitializeInternal）TpuHal/HardwareImpl19 / 0x98，20 / 0xa0按代际 HardwareImpl
TPU codec factory0x1e835fa0（TpuCodec::Create）TpuVersion switchn/a（factory）6 个按代际 CreateTpuCodec<X>
TF op-kernel 分派0xe99b000（Device::Compute）OpKernel2 / 0x10（Compute）每个 TF op kernel
MLIR pattern match0x1c9971e0（PatternApplicator::matchAndRewrite）function_ref（call *%reg）n/a（fn-ptr）每个 rewrite pattern
gRPC service handler0xf993000（RpcMethodHandler::RunHandler）std::function/AnyInvocable0x18 invoke（call *%reg）按已注册 RPC
PJRT C 入口表面0xe6aa440（GetTpuPjrtApi）扁平 PJRT_Api 结构体C fn-ptr 表每个 API 槽位一个 C callable

注意 — “Confidence CERTAIN” 表示调用点地址、偏移和槽位标签都直接读取自具名驱动函数的 IDA 反编译。两个寄存器持有目标的 HIGH 行（PatternApplicator、gRPC）对形态是确定的，但具体被调用者是动态加载的，因此无法仅从站点钉住实现者。

两种半分派形态 ​

在任何枢纽之前，先学会读取这些形态。这个二进制中的每个间接调用都是三种形式之一。

形态 A — vtable 槽位分派（多数） ​

asm
mov    (%rdi), %rax        ; load vptr from object+0  (the *(_QWORD *)obj in pseudocode)
call   *0x28(%rax)         ; call slot 5  (0x28 / 8 = 5)
```text

在反编译伪代码中，这是 `(*(...)(*(_QWORD *)obj + 0x28LL))(obj, ...)`。对象的前 8 字节是 vptr；`+0xNN` 选择槽位。这是规范的 C++ virtual call，也是下面每个 pass/visitor/kernel 枢纽所使用的方式。槽位索引是 `0xNN / 8`。

### 形态 B — 函数指针分派（`call *%reg`）

```asm
mov    0x18(%r13), %rax    ; load a function pointer out of a callable object
call   *%rax               ; no fixed offset on the call itself

伪代码是 (*(...)(a1 + 0x18))(...)，其中加载的值是原始代码指针，而不是 vptr。这是类型擦除 callable 的分派方式：llvm::function_ref<>::callback_fn、std::function、AnyInvocable，以及 mlir::PatternApplicator 的 matched-pattern predicate。grep call *0xNN(%reg) 找不到这些调用。 要枚举它们，还要扫描 call *%reg。

形态 C — 扁平 C 函数指针结构体（PJRT_Api） ​

PJRT plugin ABI 是一个包含约 140 个函数指针的 C 结构体。GetTpuPjrtApi（0xe6aa440）会构建它一次（由 __cxa_guard_acquire 保护）并返回静态实例；随后框架通过读取固定结构体偏移来调用 api->PJRT_Client_Create(...) 等。这不是 C++ vtable — 没有 this-as-first-arg 约定，也没有 RTTI 绑定。它要通过读取结构体初始化器来解析，其中每个槽位都被赋值为一个具名 TPU_PJRT_* thunk（例如 TPU_PJRT_HostAllocator_Allocate）。把它视为边界：其上是 C ABI，其下是本页其余部分映射的 C++ vtable 世界。

陷阱 — 全二进制最热的偏移是 0x10（槽位 2），但槽位 2 在每个层级中表示不同方法 — pass 的 name()、op kernel 的 Compute()、MLIR op 的 foldHook、codec 的 Encode。仅有偏移无法告诉你任何信息；必须知道对象属于哪个 vtable，槽位标签才有意义。这个连接正是 RTTI 普查提供的内容。

HLO 访问者分派 — 132 路 opcode 扇出 ​

目的 ​

XLA 层中最繁忙的单个多态分派。每个分析和优化 pass 都会遍历 HLO 图，并且在每个节点处，HloInstruction::Visit 会分派到 visitor 的逐 opcode handler。这是经典 GoF visitor：instruction 选择运行 visitor 的哪个 Handle<Opcode> 方法。

入口点 ​

text
HloInstruction::Accept (0x1e584660)        ── public visitor entry; runs the DFS, then finish-hook
  └─ PostOrderDFS<DfsHloVisitorBase> (0x1e5866e0)   ── per-node driver loop
       ├─ *0x458 ShouldProcessNode (slot 139)   ── per-node gate
       ├─ *0x448 Preprocess        (slot 137)   ── per-node pre-hook
       ├─ HloInstruction::Visit (0x1e585660)    ── the opcode fan-out (below)
       └─ *0x450 Postprocess       (slot 138)   ── per-node post-hook
  └─ *0x440 FinishVisit (slot 136)              ── post-traversal hook
```text

### 算法

```c
function HloInstruction_Visit(instr, visitor):     // 0x1e585660
    opcode = instr.byte[0xc]                        // HloOpcode at instr+12
    switch (opcode):                                // 132 cases, 0x00 .. 0x83
        case 0x6c:  return (*visitor.vtable[0x28])(visitor, instr)  // off 0x28 = slot 5
        case 0x1f:  return (*visitor.vtable[0x20])(visitor, instr)  // off 0x20 = slot 4 (lowest)
        case 0x05:  return (*visitor.vtable[0x438])(visitor, instr) // off 0x438 (highest)
        ...                                         // 132 distinct slots, one per opcode
        default:
            // "Unhandled HloOpcode for DfsHloVisitor: %s ... please file a bug for XLA."
            return Internal(...)                    // status error, opcode out of range

编译器把这个 switch 降低成 .rodata i32 jump table：movzbl 0xc(%rdi),%ecx ; cmp $0x83,%rcx ; ja default ; movslq (table,%rcx,4),%rcx ; ...; jmp *0xNN(%rcx)。每个 opcode 都尾跳到 visitor vtable 的一个不同槽位；这些偏移并不连续（从 opcode 0x1f 的 0x20 到 opcode 0x05 的 0x438），因为槽位布局遵循 DfsHloVisitor 方法声明顺序，而不是 opcode enum 顺序。

解析槽位 ​

visitor vtable 是 _ZTVN3xla17DfsHloVisitorBaseIPNS_14HloInstructionEEE @ 0x21d2c320（address point 0x21d2c330）。它的 132 个 Handle<Op> 槽位分成两类：

  • 约 56 个槽位在 base 中有默认主体 — elementwise opcode（Add、Multiply、Compare、Convert、Maximum、Negate、Tanh、…）转发到 HandleElementwiseUnary/HandleElementwiseBinary。不覆盖它们的 visitor 仍能工作。
  • 约 76 个槽位在 base 中是 __cxa_pure_virtual — 没有通用默认值的结构性 opcode（Convolution、Fusion、Dot、Reduce、CustomCall、collective）。每个具体 visitor 都必须实现这些槽位，否则无法链接。
轴值来源
Opcode0x00-0x83（132）instr+0xc 字节；Visit 中的 switch
槽位偏移0x20-0x438，不连续Visit 中的 132 个 call/jmp *0xNN 操作数
默认 vs pure约 56 个默认转发，约 76 个 purebase vtable 0x21d2c320 槽位 addend
Hook 槽位136 FinishVisit，137 Preprocess，138 Postprocess，139 ShouldProcessNodeAccept/PostOrderDFS 分派站点

怪异点 — opcode -> 槽位排列不能从 opcode enum 恢复。Visit switch 是唯一权威：opcode 0x1f 使用最低槽位（0x20），opcode 0x05 使用最高槽位（0x438）。要重建 visitor 契约，必须读取 132 个 case 主体，而不是假设槽位顺序匹配 enum 顺序。精确的 enum-name-to-slot 映射需要解码 .rodata jump table 并连接到 HloOpcode enum；这里不复现该连接。

HLO Pass 分派 — 六字节 trampoline ​

目的 ​

HLO pipeline 中的每个 pass 都汇入两个六字节尾跳之一。HloPassInterface::Run 不是带主体的方法，而是加载 pass 的 vtable 并跳到逐 pass 实现槽位的 trampoline。

算法 ​

c
function HloPassInterface_Run(pass, module_and_set):    // 0x1e472a60, 6 bytes
    return (*pass.vtable[0x28])()                         // jmp *0x28(%rax) = slot 5 = RunImpl

function HloPassInterface_Run_uptr(pass, uptr_and_set):  // 0x1e472a80, 6 bytes
    return (*pass.vtable[0x30])()                         // jmp *0x30(%rax) = slot 6 = RunImpl(uptr)
```text

pipeline driver `HloPassPipeline::RunPassesInternal<HloModule*>`（`0x1c83ddc0`）并不内联调用槽位 5。它分派外围元数据槽位，然后调用 `pass->Run()`（非 virtual trampoline），后者尾跳到真正主体。driver 中观察到的元数据分派：

| Off | 槽位 | 方法 | 用途 |
|---|---|---|---|
| `0x10` | 2 | `name()` | logging / dump（driver 中调用 ×6） |
| `0x18` | 3 | `RunOnChangedComputations` | 每个 pass 一次 |
| `0x20` | 4 | `IsPassPipeline()` | 每个 pass 一次（×2） |
| `0x30` | 6 | `RunImpl(uptr&)` | 嵌套 pipeline 路径（×2） |

> **注意 —** 因为主体是通过 `Run` trampoline 到达，而不是内联到达，所以停在 `RunPassesInternal` 的静态 caller-graph 会完全漏掉逐 pass 工作。跟随 `0x1e472a60` 处的 trampoline 找到槽位 5，然后通过 RTTI 普查枚举槽位 5 的实现者，以列出每个具体 pass。

---

## MLIR Pass 分派 — `OpToOpPassAdaptor::run`

### 目的

MLIR pass 基础设施基于 CRTP；adaptor `mlir::detail::OpToOpPassAdaptor::run`（`0x1cb6dc20`）是驱动器，会为每个被调度的嵌套 operation 调用一次具体 pass 的 `runOnOperation()`。

### 算法

```c
function OpToOpPassAdaptor_run(pass, op, am, ...):   // 0x1cb6dc20
    if op.name.impl.typeid == UnregisteredOp:
        emitOpError("trying to schedule a pass on an unregistered operation")
        return failure
    ...
    (*pass.vtable[0x10])(pass)                        // call *0x10(%rax) = slot 2 = getName()  (logging)
    (*pass.vtable[0x20])(pass, IsIsolatedFromAbove)   // call *0x20(%rax) = slot 4 = hasTrait query
    (*pass.vtable[0x50])(pass, op)                    // call *0x50(%rax) = slot 10 = canScheduleOn
    ...
    (*pass.vtable[0x38])(pass)                        // call *0x38(%rax) = slot 7 = runOnOperation

adaptor 内读取 pass vtable 的分派站点（mov (pass),%rax ; call *0xNN(%rax)）：*0x10（槽位 2，getName）、*0x20（槽位 4，门控调度的 hasTrait<IsIsolatedFromAbove> 查询）、*0x50（槽位 10，canScheduleOn(Operation*)）以及 *0x38（槽位 7，runOnOperation — 逐 pass 主体，经由 runOnOperationImpl/runOnOperationAsyncImpl 到达）。槽位 7 是唯一运行用户 pass 逻辑的槽位；其余是自动生成的 *PassBase CRTP 元数据。

Off槽位方法
0x387runOnOperation() — pass 主体
0x102getName()
0x204hasTrait<IsIsolatedFromAbove>() 查询
0x5010canScheduleOn(Operation*)

注意 — driver 还包含 call *0x20、*0x28 和 *0x30 站点，它们分派的是另一个对象 — 它所迭代的 PassInstrumentation 列表（mov (list[i]),%rax ; call *0xNN(%rax)），也就是逐 pass 的 runBeforePass/runAfterPass/runAfterPassFailed 回调，而不是 pass vtable。把这个函数中的每个间接调用都归因于 pass vtable 的扫描会误标这些 instrumentation 槽位；只有上面的四个站点读取 pass 对象本身。

MLIR Op-Model 分派 — Operation::fold 和 concept 对象 ​

目的 ​

MLIR 不会为每种 operation 类型保留一个 C++ vtable。每个已注册 op 都有一个 OperationName::Impl — 一个“Model” concept 对象，其 vtable 携带该 op 的 hook（foldHook、hasTrait、getCanonicalizationPatterns、…）。mlir::Operation::fold（0x1d8cd480）是所有已注册 op 的中心 folding 入口；它通过这个 concept 间接层到达 op 行为，而不是通过 Operation 对象自身的 vtable。

算法 ​

c
function Operation_fold(op, attrs, results):       // 0x1d8cd480
    model = op.field[0x30]                           // mov 0x30(%rdi),%rdi  — OperationName::Impl
    ok    = (*model.vtable[0x10])(model, op, ...)    // call *0x10(%rax) = slot 2 = foldHook
    if ok: return true
    // fallback: look up a DialectFoldInterface and try its fold
    dialect = model.dialect
    if isa<DialectFoldInterface>(dialect):
        iface = dialect.interface_map[TypeID(DialectFoldInterface)]
        if iface: return (*iface.vtable[0x10])(iface, op, ...)   // a second slot-2 dispatch
    return false
```text

加载链证明这是一个真实的 Model 分派，而不是 member-pointer：`mov 0x30(%rdi),%rdi`（加载 Model concept）、`mov (%rdi),%rax`（加载其 vptr）、`call *0x10(%rax)`（槽位 2 = `foldHook`）。同一个 concept 对象还通过相邻槽位回答 `hasTrait`。`OperationName::Impl` 是逐 op 生成的分派对象，用来代替 MLIR 有意不发射的逐 op vtable。

> **怪异点 —** `Operation::fold` 中有**两个** `call *0x10` 分派：先是 op 自身的 `foldHook`，然后 — 如果它返回 false — 是 op 所属 dialect 的 `DialectFoldInterface::fold`。两者都是槽位 2 调用，但经由不同 concept 对象（op Model vs dialect interface）。停在第一个 `call *0x10` 的导航者会漏掉 dialect 级 fold fallback。

---

## CPU Thunk 执行 — `Thunk::Execute` = 槽位 5

### 目的

XLA:CPU backend 把 computation 降低成一系列 `Thunk` 对象；`ThunkExecutor::TracedExecute`（`0x1c0f0320`）通过分派其 `Execute` 主体来运行一个 thunk。

### 算法

```c
function ThunkExecutor_TracedExecute(executor, thunk, params):   // 0x1c0f0320
    if g_trace_level > 0:
        TraceMeProducer(...)                                       // profiler scope
        (*thunk.vtable[0x28])(executor, thunk, params)             // call *0x28 = slot 5 = Execute
        AsyncValueRef<Chain>::AndThen(...)                         // chain continuation
        TraceMe::Stop(...)
    else:
        (*thunk.vtable[0x28])(executor)                            // same slot 5, untraced fast path

traced 和 untraced 路径都分派同一槽位 — 0x28 = 槽位 5 = xla::cpu::Thunk::Execute(ExecuteParams const&)。ThunkExecutor::ExecuteSequential 对序列中的每个 thunk 调用一次 TracedExecute；async continuation 路径馈入同一个槽位 5 站点。CpuExecutable::ExecuteThunks 是构建 executor 的顶层入口。

TPU ISA Codegen — EmitInstruction 跨 IsaEmitter 扇出 ​

目的 ​

xla::jellyfish::CodeGenerator::EmitInstruction（0x14043a40）是逐 LloInstruction codegen 分派器。它是最密集的单个多态分派区域：一个大函数（3,652 行反编译代码）在 152 槽 IsaEmitter vtable 中跨 81 个不同槽位扇出。具体的逐代际 emitter 会填充这些槽位。

算法 ​

c
function EmitInstruction(codegen, llo_instr, bundle):   // 0x14043a40
    emitter = codegen.isa_emitter                         // the per-gen IsaEmitter object
    switch (llo_instr.kind):                               // 81 reachable emitter slots
        ...
        case VectorMatmulMsk:
            (*emitter.vtable[0x418])(emitter, ...)          // off 0x418 = slot 131 = EmitVectorMatmulMsk
        case AccumulatorBinop:
            (*emitter.vtable[0x478])(emitter, ...)          // off 0x478 = slot 143 = EmitVectorAccumulatorBinop
        ...                                                 // EmitVectorMove/Pack, transpose, Cmem,
                                                            // BarnaCore sync/wait, event/program hooks
```text

该分派是真正经由 emitter vptr 的 virtual dispatch：`(*(...)(*(_QWORD *)emitter_obj + 0x418LL))(IsaEmitter*, ...)`。已确认的站点包括 `0x418`（槽位 131，`EmitVectorMatmulMsk`）和 `0x478`（槽位 143，`EmitVectorAccumulatorBinop`）。这些槽位是由 `{Pf, Vf, Gl, Gf}` 具体 emitter 类填充的逐代际 ISA 编码器 hook；有关如何选择每个代际的 emitter，请参见[逐代际函数分派](per-gen-function-dispatcher.md)。

> **注意：** 扇出是 **81** 个不同的 `IsaEmitter` vtable 槽位偏移，跨越 `0x50`-`0x490`。该计数是 `0x14043a40` 反编译中接收者被 IDA 类型化为 `xla::jellyfish::IsaEmitter *` 的不同分派操作数数量；同一函数中少数其他间接调用以 `LloInstruction`/helper 对象为目标，而不是 emitter，因此被排除。

---

## TpuHal 硬件启动 — 槽位 19 和 20

### 目的

`tpu::TpuHal::InitializeInternal`（`0x1e811ea0`）是硬件启动驱动器。真正的逐代际多态集中在两个 `HardwareImpl` 槽位；周围大多数 `TpuHal::` 表面都是对 chip vector 和 memory manager 的非 virtual 委托。

### 算法

```c
function TpuHal_InitializeInternal(hal, options):     // 0x1e811ea0
    validate(options)                                   // page-alignment / power-of-2 checks
    InitializeAllocator(hal, ...)                       // non-virtual
    status = (*hal.vtable[0x98])(hal)                   // slot 19 — per-gen ValidateTopology
    if status != ok: return status
    TpuHalCommonStates::Create(...)
    (*hal.vtable[0xa0])(&out, hal, options)             // slot 20 — CreateAndInitializeChips (per-gen)
    sort(chips); for each chip: (*hal.vtable[0x48])(hal, i)   // slot 9 — GetChip (per-chip)
    ParallelForWithStatus(...)                          // bring up chips in parallel
    on error: (*hal.vtable[0x20])(hal)                  // slot 4 — TearDown

槽位 19（0x98）是 topology validation hook，槽位 20（0xa0）是 CreateAndInitializeChips — 二者都由逐代际 HardwareImpl 子类覆盖。driver 还会按 chip 分派槽位 9（0x48，GetChip），并在 error path 上分派槽位 4（0x20，TearDown）。

注意： GetChip 在 InitializeInternal 的逐 chip validation loop 中通过槽位 9（call *0x48(%rax)）被虚分派，而不是作为非 virtual chip-vector 访问。其他具名 TpuHal:: 方法（InitializeAllocator、page-alignment checks）是非 virtual；GetChip 不是。

TPU Codec Factory — 一个 TpuVersion switch，然后是 6 槽 vtable ​

目的 ​

tpu::TpuCodec::Create（0x1e835fa0）是一个逐代际 factory — 一个 switch(TpuVersion)，不是 virtual call。它返回一个 codec 对象，随后该对象携带 6 槽 TpuCodec vtable，消费者通过它调用 Encode/Decode。

算法 ​

c
function TpuCodec_Create(out, version):       // 0x1e835fa0
    switch (version):
        case 0: codec = CreateTpuCodecJellyfish(out)
        case 1: codec = CreateTpuCodecDragonfish(out)
        case 2: codec = CreateTpuCodecPufferfish(out)
        case 3: codec = CreateTpuCodecViperfish(out)
        case 4: codec = CreateTpuCodecGhostlite(out)
        case 5: codec = sub_1E838380(out)         // anonymous v5 codec
    out.tag   = 1                                 // variant discriminant
    out.codec = codec                             // out+8 = the codec object
    return out
```text

返回对象暴露一个 6 槽 `TpuCodec` vtable：`Encode`（槽位 2 / `0x10`）、`Decode`（槽位 3 / `0x18`）、`EncodeBundle`（槽位 4 / `0x20`）、`DecodeBundle`（槽位 5 / `0x28`）。encode/decode 消费者经由 `pro::v4::proxy` facade 路由 — 这是第三种分派机制（proxy-table），不同于 C++ vtable 和函数指针；其内部槽位布局未在此解码。

> **注意：** `TpuVersion` case-to-codec 映射是 **`{0:Jellyfish, 1:Dragonfish, 2:Pufferfish, 3:Viperfish, 4:Ghostlite, 5:anon-v5}`**，直接读取自 `0x1e835fa0` 反编译 — case 1 是 Dragonfish，case 4 是 Ghostlite。这与[逐代际分派器](per-gen-function-dispatcher.md)页面上的 `TpuVersion` ordinal ladder 匹配。

---

## TensorFlow Op-Kernel 分派 — `Device::Compute` = 槽位 2

### 目的

`tensorflow::Device::Compute`（`0xe99b000`，12 字节）是规范的 TF op-kernel 分派：device 通过尾跳到 kernel 的 `Compute` 槽位来运行 kernel。

### 算法

```c
function Device_Compute(device, kernel, ctx):    // 0xe99b000
    return (*kernel.vtable[0x10])(kernel, ctx)     // jmp *0x10 = slot 2 = OpKernel::Compute

偏移 0x10 = 槽位 2 = OpKernel::Compute(OpKernelContext*)。ThreadPoolDevice::Compute（0x10835420）执行同样的 call *0x10；RenamedDevice::Compute（0x108a5200）则通过 jmp *0xb8(%rax) 委托给被包装的 device。每次 op 执行触发一次。

函数指针和 C-ABI 枢纽 ​

这些枢纽不使用形态 A。RTTI 无法解析它们；需要读取调用者来找出寄存器/结构体持有什么。

MLIR pattern match — call *%reg ​

mlir::PatternApplicator::matchAndRewrite（0x1c9971e0）应用 rewrite pattern。它的 match predicate 是作为参数传入的 llvm::function_ref，并以 a4(arg, pattern) 调用 — 一个裸 call *%reg，不是内联 vtable 槽位。匹配到的 pattern 自身的 matchAndRewrite 同样通过从 pattern 对象加载的寄存器持有指针到达。槽位在二进制中不是固定的：它从每个 Pattern 动态加载，因此没有单个偏移可锚定。

gRPC service handler — call *0x18(%reg) ​

模板化的 grpc::internal::RpcMethodHandler<...>::RunHandler（0xf993000）调用已注册的 service-method callable：(*(...)(handler + 0x18))(...)。偏移 0x18 选择类型擦除 callable（std::function/AnyInvocable）的 invoke 指针，这同样是函数指针形态。具体 service method 是 service-build 时注册的任何目标。

PJRT C 入口表面 — 扁平结构体 ​

pjrt::tpu_plugin::GetTpuPjrtApi（0xe6aa440）返回静态 PJRT_Api 结构体，并在 __cxa_guard 下惰性构建它的 extension chain（raw-buffer、layouts、memory-descriptions、executable-metadata、host-allocator、cross-host-transfers）。框架通过固定结构体偏移调用 API 槽位；每个槽位持有一个具名 TPU_PJRT_* C thunk。这是分派栈顶部 — 解析 PJRT 调用时先读取结构体初始化器，然后跟随具名 thunk 进入 C++ 世界。

陷阱 — 函数指针枢纽对 call *0xNN(%reg) 扫描是不可见的。要枚举类型擦除和 pattern-rewrite 分派，必须也扫描 call *%reg。跳过这一步会让整个 MLIR pattern-rewrite 和 gRPC service 表面保持未映射。

解析间接调用 — 方法论 ​

这些枢纽展示的流程如下。给定一个想要跟踪的未知 call *0xNN(%reg)：

text
1. CLASSIFY THE SHAPE
   a. backtrack the register: is it `mov (obj),%vptr` (Shape A) or a loaded code ptr (Shape B)?
   b. if the callee came from a flat struct read by fixed offset -> Shape C (C-ABI), stop here,
      read the struct initializer.

2. IDENTIFY THE VTABLE  (Shape A only)
   a. find where `obj` was constructed or typed (ctor call, factory return, RTTI check).
   b. the object's class names the vtable: look it up in the RTTI census by mangled name.
      e.g. an object whose vptr is 0x21d2c330 -> DfsHloVisitorBase<HloInstruction*> @ 0x21d2c320.

3. COMPUTE THE SLOT
   slot = 0xNN / 8.   (0x28 -> slot 5, 0x10 -> slot 2, 0x38 -> slot 7.)

4. LABEL THE SLOT
   the RTTI census gives the per-slot method name for that vtable
   (slot 5 of HloPassInterface = RunImpl; slot 2 of OpKernel = Compute; slot 7 of mlir::Pass
    = runOnOperation).

5. ENUMERATE IMPLEMENTORS
   read the slot addend of EVERY vtable bound to that base class in the RTTI census.
   each distinct addend at offset 0xNN is one candidate concrete target.
   the *set* of those addends is the fan-out of the call site.

6. PRUNE BY CONTEXT
   the construction site (step 2) usually fixes which concrete subclass `obj` is, collapsing
   the fan-out to the one or few implementors actually reachable from this caller.
```text

对于形态 B（`call *%reg`），步骤 2-5 不适用：没有 vptr，也没有固定槽位。改为把寄存器回溯到其定义 — 一个 `callback_fn` 实例化、一个 `std::function` 目标赋值，或一次 `Pattern` 字段加载 — 并在那里读取捕获的目标。对于形态 C，实现者在结构体初始化器中直接具名。

> **注意 —** 槽位索引 /8 规则和 RTTI 普查连接就是全部关键。本页上的每个枢纽都正是通过步骤 1-5 解析的：读取 driver，计算 `0xNN/8`，在普查中查找槽位，枚举 addend。构建导航器的重新实现者应当自动化这个连接，而不是手工解析每个站点。

---

## 相关组件

| 组件 | 关系 |
|---|---|
| [分派表分类](dispatch-table-taxonomy.md) | 分类每个枢纽读取的是*哪种*表（vtable、factory switch、C-ABI struct、proxy facade） |
| [RTTI / Vtable 普查](rtti-vtable-census.md) | 被调用方侧：逐槽位方法标签，以及本页连接到的实现者集合 |
| [逐代际函数分派](per-gen-function-dispatcher.md) | `TpuCodec::Create` / `TpuHal` 槽位 20 如何选择逐代际实现者 |

## 交叉引用

- [取证概览](overview.md) — 多态分派在二进制解剖图中的位置
- [分派表分类](dispatch-table-taxonomy.md) — 枚举表*种类*的姊妹篇；本页枚举进入这些表的*调用点*
- [RTTI / Vtable 普查](rtti-vtable-census.md) — 实现者查找；上面的解析步骤 4-5 读取其逐槽位标签和 addend
- [逐代际函数分派](per-gen-function-dispatcher.md) — `IsaEmitter` / TpuHal / TpuCodec 枢纽背后的 `{Pf,Vf,Gl,Gf}` emitter 和 codec/HAL 逐代际选择