TpuCompiler 名册

本页所有地址均适用于来自 libtpu-0.0.40-cp314 wheel 的 libtpu.so（构建 libtpu_lts_20260413_b_RC00，build-id md5 89edbbe81c5b328a958fe628a9f2207d，ELF x86-64 DYN，未 stripped；反修饰后的 C++ 名称和 IDA 恢复的 C 名称均按原文引用）。.text VMA 等于文件偏移。其他版本会不同。

摘要

TpuCompiler_* 是 StreamExecutor TPU 编译器的 C-ABI 正面接口：一个扁平的 extern "C" 表面，开源 xla::TpuCompiler shim（构建进 TensorFlow/XLA，不构建进此二进制）通过它驱动闭源 libtpu.so 编译流水线，而无需在 .so 边界共享任何 C++ 类型。宿主侧 shim 具有名为 RunHloPasses、RunBackend、Compile、ShapeSize 和 DefaultDeviceShapeRepresentation 的方法；每个方法都会通过 TfTpu_ExecutorApiFn 的一个函数指针槽转发到这里列出的七个 TpuCompiler_* 自由函数之一。它们每一个都只接收不透明句柄、一个序列化 proto blob 和一个 status-out 对象，在 libtpu 自身静态链接的 XLA 内部重建丰富的 xla:: 对象，调用真实的编译器 vtable，然后重新序列化结果。句柄背后的编译器对象是 xla::jellyfish::DeepseaCompiler，即 xla::Compiler 的 TPU（“Deepsea”/Jellyfish）子类。

第二个在词法上相邻的集群 TpuCompile_*（六个函数）不是同一个表面，并且经常造成混淆。TpuCompiler_* 是通过 ExecutorApiFn 到达的 SE 类方法 C-API；TpuCompile_* 是 TensorFlow TPUCompileOp kernel 的独立支持层：一次性 CompileAndBuild 会端到端运行 tensorflow::tpu::TpuCompileOpKernelCommon 并发出 XLA_TpuProgram 句柄，此外还包含编译缓存键构造、guaranteed-const 指纹和两个策略谓词。这两个集群共享 tpu_executor_c_api.cc / tpu_util_c_api.cc 翻译单元家族和 Compile 这个动词，但它们处于不同层级：TpuCompiler_* 是编译器类的一个方法，TpuCompile_CompileAndBuild 是整个 TF compile op。

本页负责两个集群的逐函数名册和 impl-symbol 映射。ABI 接缝本身，即为什么存在扁平 C 接口、*ApiFn() 访问器模式以及不透明句柄 / ApiConverter::ToC/FromC 约定，已在 shim 概览中一次性建立，这里不再重述。这些函数调用的 HLO pass 调度由 HLO Pass Registry负责；它们产生的 XLA_TpuProgram 句柄由 TpuProgram 名册负责。

对于重新实现，契约是：

七个 TpuCompiler_* 签名和分发：每个签名跨越哪种 proto（HloModuleProto、HloModuleGroupProto）、每个索引哪个编译器 vtable 槽（+24 RunHloPasses、+32 RunBackend、+96 ShapeSize、+104 DefaultDeviceShapeRepresentation），以及句柄的双指针形态。
proto-in / proto-out 编组纪律：输入经 DeserializeProto，经 CreateFromProto 构建 xla::HloModule，运行，再用 ToProto + SerializePartialToArray 输出到调用方释放的字节缓冲区。
DeepseaCompiler 对象模型：句柄是 Compiler**；真实对象是 xla::jellyfish::DeepseaCompiler；RegisterAllPhases 是一次性的 phase 注册钩子，由开源 TpuCompiler 构造函数调用（“Initialize”步骤）。
TpuCompile_* 支持集群：CompileAndBuild 运行什么、缓存键/指纹辅助函数，以及中止抑制策略谓词。


句柄类型	`xla::jellyfish::DeepseaCompiler**`（双指针；外层 = 8 字节盒子，内层 = 编译器）
构造函数	`TpuCompiler_New @ 0xeabc4a0` → `operator new(8)` ×2，`DeepseaCompiler::DeepseaCompiler`
析构函数	`TpuCompiler_Free @ 0xeabc4e0` → 虚析构槽 `+8`，然后 `free` 盒子
Compile 分发	`xla::Compiler` vtable：RunHloPasses `+24`、RunBackend `+32`、Compile 经由 `xla::Compiler::Compile`、ShapeSize `+96`、DefaultDeviceShapeRepresentation `+104`
输入线格式	`TpuSerializedProto`（ptr+len）→ `stream_executor::tpu::DeserializeProto<…>`
状态通道	尾部 `absl::status_internal::StatusRep**` out-param（SE “ok?-code-msg” 惯用法）
C-ABI 源 TU	`learning/45eac/tfrc/executor/stream_executor/tpu_executor_c_api.cc`（compiler），`tpu_util_c_api.cc`（TpuCompile helpers）
Phase 注册钩子	`xla::TpuCompiler::RegisterAllPhases() @ 0xf849ec0`
名册规模	7 个 `TpuCompiler_` + 6 个 `TpuCompile_` = 13 个自由函数
证据等级	重新实现级 / 已按 IDA 反编译逐字节确认

注意 — 此构建中没有 TpuCompiler_Initialize 符号。重新实现者预期的 “initialize” 步骤被拆成两部分：对象构造是 TpuCompiler_New，一次性 phase 表注册是 xla::TpuCompiler::RegisterAllPhases（0xf849ec0），它由宿主侧 xla::TpuCompiler 构造函数调用，而不是由 C-ABI 调用。同样也没有 TpuCompiler_*ToTpuProgram；产生序列化 XLA_TpuProgram 是第二个集群中 TpuCompile_CompileAndBuild 的工作，而 TpuCompiler_RunBackend 返回进程内 xla::Executable* 句柄，不返回序列化程序。

1. 生命周期 — `New` / `Free`

目的

构造和销毁不透明编译器句柄。宿主侧 xla::TpuCompiler shim 在其生命周期内持有 New 返回的句柄，并通过 Free 释放它。编译器跨调用无状态：RunHloPasses / RunBackend / Compile 每次都接收句柄和一个新的 module，因此单个句柄服务整个进程。

算法

// TpuCompiler_New                                              0xeabc4a0
DeepseaCompiler** TpuCompiler_New():
    box   = operator new(8)                  // the handle the host receives
    inner = operator new(8)                  // the actual compiler object (8-byte: just a vptr)
    DeepseaCompiler::DeepseaCompiler(inner)  // base xla::Compiler ctor; sets vtable
    *box = inner
    return box                               // host holds a Compiler**

// TpuCompiler_Free                                             0xeabc4e0
void TpuCompiler_Free(Compiler** box):
    if *box:
        (*box)->vtable[+8](*box)             // virtual destructor — DeepseaCompiler::~DeepseaCompiler
    free(box)                                // release the outer box

怪异点 — 句柄是一个双重指针。New 分配一个 8 字节盒子，分配 8 字节编译器，并把编译器存入盒子；其他每个 TpuCompiler_* 函数都会解引用两次（**a1）以到达 vtable。重新实现者如果直接返回编译器指针，会让所有后续 (*(_QWORD *)*a1 + off) vtable 索引错位。编译器对象本身只有 8 字节：它是 vtable 之上的纯虚行为，没有实例状态；因此 DeepseaCompiler 通过每次调用传入的 CompileOptions 携带配置，而不是通过字段携带。

函数映射

函数	地址	角色
`TpuCompiler_New`	`0xeabc4a0`	分配盒子 + `DeepseaCompiler`，返回 `Compiler**`
`TpuCompiler_Free`	`0xeabc4e0`	虚析构编译器，释放盒子

注意事项

New 不接收参数，并且在观测到的函数体中不会失败（没有 status-out；原始 operator new 在 OOM 时会 std::terminate，而不是返回 null）。Free 容忍空内层指针（if (*ptr)），但不容忍空盒子。phase 表注册，即重新实现者可能以为在 New 内部发生的部分，发生在别处：见 §5。

2. HLO Passes 与 Backend — `RunHloPasses` / `RunBackend`

目的

XLA 两阶段编译的两个半部。RunHloPasses 在一个 module 上运行目标无关再到 TPU 的 HLO 优化流水线，并返回优化后的 module（重新序列化）。RunBackend 消费已经优化的 module，并将其 lower 成后端 xla::Executable，作为进程内句柄返回。当宿主 shim 希望在两个阶段之间插入处理（例如缓存优化后的 HLO）时，会连续调用它们；Compile（§3）是融合形式。

入口点

text

xla::TpuCompiler::RunHloPasses(...)            (host-side shim, not in this binary)
  └─ ExecutorApiFn()->slot[RunHloPasses]       0x20819360 = ExecutorApiFn() accessor
       (returns &executor_api_fn @ 0x2258c818; the RunHloPasses fn-ptr lives in that struct)
       └─ TpuCompiler_RunHloPasses             0xeabcd80   — C-ABI impl
            └─ DeepseaCompiler vtable[+24]      xla::Compiler::RunHloPasses override

RunBackend 与之相同，只是槽/偏移为 +32。

算法

// TpuCompiler_RunHloPasses                                     0xeabcd80
// args: a1=Compiler** handle, a2=TpuSerializedProto* (HloModuleProto in),
//       a3=XLA_HloModuleConfig*, a4=SE_StreamExecutor* (may be null),
//       a5=TpuSerializedProto* (out), a6=StatusRep** (status out)
int TpuCompiler_RunHloPasses(a1, a2, cfg, exec, out, status):
    proto = DeserializeProto<HloModuleProto, TpuSerializedProto>(a2)   // ptr+len → proto
    config = ApiConverter::FromC(cfg)                                  // XLA_HloModuleConfig → xla::HloModuleConfig
    module = HloModule::CreateFromProto(proto, config, /*flags*/1,0,1,0)
    if module is error: write status; return                          // StatusOr unwrap

    // build a CompileOptions whose device allocator wraps the SE executor, if given
    opts = {}
    if exec && exec->allocator:
        opts.device_allocator = new WrapperDeviceMemoryAllocator{      // vtable off_21616EF8
            platform = GetUnderlyingDeepseaPlatform(),                 // Meyers singleton
            executor = exec }
    opts.layout_canonicalization_callback = empty                      // default-constructed std::function

    optimized = (**a1).vtable[+24](compiler, module, exec, opts)       // xla::Compiler::RunHloPasses
    if optimized is error: write status; goto cleanup

    // re-serialize the optimized module back across the seam
    p = HloModuleProto(); optimized->ToProto(&p)
    n = p.ByteSizeLong()
    buf = operator new(n)
    if !p.SerializePartialToArray(buf, n): LOG(FATAL) "proto_helper.h:45"
    out->ptr = buf; out->len = n                                       // caller frees buf
cleanup:
    destroy module/optimized/opts; return

RunBackend（0xeabd100）在分发前结构相同，随后在结果上分叉：

// TpuCompiler_RunBackend                                       0xeabd100  (vtable +32)
    executable = (**a1).vtable[+32](compiler, module, exec, opts)  // xla::Compiler::RunBackend → xla::Executable*
    if ok:
        h = operator new(8); *h = executable                       // box the Executable*
        *out = h                                                   // out is an SE_Executable** handle, NOT a proto

怪异点 — RunHloPasses 返回重新序列化的 HloModuleProto（调用方释放的字节缓冲区），但 RunBackend 返回一个装入 8 字节分配中的活 xla::Executable*；后端结果从不以 proto 形式跨越接缝。这种不对称是有意的：优化后的 HLO 可移植且可缓存，因此会序列化；Executable 持有设备驻留状态，并在进程内消费，因此只回传其指针（一个不透明 TpuExecutable 句柄）。
陷阱 — 两个函数都接收一个可选 SE_StreamExecutor*（a4）。当它非空时，impl 会惰性构造一个 WrapperDeviceMemoryAllocator（vtable off_21616EF8，在 stream_executor::DeviceMemoryAllocator 上实现 Allocate/Deallocate/GetStream），绑定到进程级 DeepseaPlatform 单例（GetUnderlyingDeepseaPlatform，由 _cxa_guard 保护），并把它穿入 CompileOptions.device_allocator。忽略 a4 的重新实现会在没有 device allocator 的情况下编译，并静默禁用 allocation-aware passes（例如需要真实 HBM 大小的 memory-space assignment）。

函数映射

函数	地址	Vtable 槽	结果形式
`TpuCompiler_RunHloPasses`	`0xeabcd80`	`+24`	重新序列化的 `HloModuleProto`（字节 buf）
`TpuCompiler_RunBackend`	`0xeabd100`	`+32`	装盒的 `xla::Executable*` 句柄

注意事项

CreateFromProto 标志 (1,0,1,0) 请求 prohibit-ill-formed / no-verifier-on-the-cheap-path 行为（精确标志语义 LOW：来自参数位置，而不是 verifier trace）。SerializePartialToArray 失败时的 LOG(FATAL)（proto_helper.h:45）意味着重新实现必须保证缓冲区在序列化前按 ByteSizeLong() 定大小：欠大的缓冲区没有错误返回，只有进程中止。

3. Compile — 融合入口 {#3-compile--the-fused-entry}

目的

TpuCompiler_Compile 是单次调用形式：它接收一个 HloModuleGroupProto，构建 module，并通过 xla::Compiler::Compile 运行完整的 optimize-then-lower 流水线，把每个 module 的一个装盒 Executable* 返回到调用方提供的数组中。这是 PJRT 适配器和大多数调用方在不需要介入 HLO passes 与 backend 之间时使用的路径。

算法

// TpuCompiler_Compile                                          0xeabc520
// a1=Compiler** handle, group proto in, a5/a3 = module count guard,
// per-executor list in, a6=array-of-StreamExecutor-lists, a7=out Executable* array, a8=StatusRep**
void TpuCompiler_Compile(a1, group, ..., out_array, status):
    grp = DeserializeProto<HloModuleGroupProto, TpuSerializedProto>(...)
    if module_count > 1:                                          // a5<=1 && v54<2 guard
        *status = MakeError("Can not compile multiple HLO modules at once.")   // c_api.cc:1040
        return
    config = ApiConverter::FromC(...)
    if grp.modules_size() <= 0: LogIndexOutOfBoundsAndAbort()      // proto bounds check
    module = HloModule::CreateFromProto(grp.modules(0), config, 1,0,1,0)
    if error: write status; cleanup

    // install the layout-canonicalization callback (TpuCompiler_Compile::$_0) on the module config
    module.config.set_layout_canonicalization_callback(&$_0)       // module+3864/+3872 fn-ptr pair

    // flatten the per-module vector<StreamExecutor*> (loop-unrolled by 8)
    execs = copy_stream_executors(...)
    opts  = CompileOptions{ device_allocator = WrapperDeviceMemoryAllocator(platform, execs) if present }
    result = xla::Compiler::Compile(compiler, module, execs, opts)  // StatusOr<vector<unique_ptr<Executable>>>
    if ok:
        for i in result:                                           // box each Executable*, transfer ownership
            h = operator new(8); *h = result[i].release()
            out_array[i] = h
    else:
        write status
    free temporaries; cleanup

函数映射

函数	地址	输入 Proto	结果形式
`TpuCompiler_Compile`	`0xeabc520`	`HloModuleGroupProto`	装盒 `Executable*` 的数组

注意事项

陷阱 — 单次调用入口接收一个 HloModuleGroupProto（一个组），但拒绝任何包含超过一个 module 的组：在 tpu_executor_c_api.cc:1040 处为 "Can not compile multiple HLO modules at once."。组形输入是通用 xla::Compiler::Compile(HloModuleGroup, ...) 签名的遗留；TPU C-API 每次调用只支持一个 module。重新实现者必须保留组容器，但强制 size-1 不变量，并且仍必须对 modules(0) 做边界检查（impl 在空组上调用 LogIndexOutOfBoundsAndAbort）。
怪异点 — Compile 在分发前直接把一个布局规范化回调（TpuCompiler_Compile::$_0）安装到 module config（module+3864/+3872）上，而 RunHloPasses/RunBackend 安装的是空回调。这就是融合路径允许后端重新规范化入口计算布局的方式，而拆分路径把这件事留给宿主。把 Compile 复现为 “RunHloPasses then RunBackend” 但不带此回调，会在布局敏感 module 上产生分歧。

每个 module 的 executor 列表通过 8 路展开循环复制，并且源指针被双重解引用（**(_QWORD**)(base + 8*i)）：由于只允许一个 module，输入是展平为一个内层列表的 vector<vector<StreamExecutor*>>。

4. Shape 查询 — `ShapeSize` / `DefaultDeviceShapeRepresentation`

目的

宿主用于缓冲区定大小和布局的两个编译器纯函数，独立于任何 compile。ShapeSize 返回一个 shape 在设备上占用的字节数；DefaultDeviceShapeRepresentation 把宿主（逻辑）shape 映射为 TPU 实际存储的设备（物理、tiled/padded）shape。

算法

// TpuCompiler_ShapeSize                                        0xeabd400
int64 TpuCompiler_ShapeSize(Compiler** a1, XLA_Shape* a2):
    shape = ApiConverter::FromC(a2)                          // XLA_Shape → xla::Shape
    fn = (**a1).vtable[+96]()                                // returns a ShapeSizeFunction object (closure)
    size = fn.call(shape)                                    // invoke the size functor
    fn.dtor()                                                // release the functor
    return size

// TpuCompiler_DefaultDeviceShapeRepresentation                0xeabd480
void TpuCompiler_DefaultDeviceShapeRepresentation(Compiler** a1, XLA_Shape* in, XLA_Shape* out):
    host_shape = ApiConverter::FromC(in)
    dev_shape  = (**a1).vtable[+104](compiler, host_shape)   // xla::Compiler::DefaultDeviceShapeRepresentation
    ApiConverter::ToC(dev_shape, out)                        // xla::Shape → XLA_Shape (caller owns 'out')

注意 — ShapeSize 不直接调用 vtable 方法；槽 +96 返回一个函子（类似 std::function 的 ShapeSizeFunction，6-qword v6 块：vtable、state、invoke-ptr、deleter），随后 C-ABI 调用并销毁它。这层间接让编译器能够把其 size 策略作为一等 callable 暴露给宿主。重新实现者必须调用返回的函子并运行其 deleter（函子 vtable 的槽 +8），不能假设 +96 本身就是 sizer。

函数映射

函数	地址	Vtable 槽	返回
`TpuCompiler_ShapeSize`	`0xeabd400`	`+96`（函子工厂）	`int64` 字节大小
`TpuCompiler_DefaultDeviceShapeRepresentation`	`0xeabd480`	`+104`	设备 `XLA_Shape`（out-param）

注意事项

DefaultDeviceShapeRepresentation 通过 ApiConverter::ToC 填充调用方提供的 XLA_Shape，因此调用方拥有结果，并且必须将其与 ApiConverter::Destroy(XLA_Shape*)（内部释放重载）配对。它返回的设备 shape 是 TPU 对输入宿主 shape 使用的 tiled/padded 表示，这是 XLA 在 libtpu 不暴露 layout assignment 内部的情况下学习某个缓冲区 TPU (8, 128) 风格 tiling 的机制。

5. Phase 注册 — “Initialize” 步骤 {#5-phase-registration-the-initialize-step}

目的

开源 xla::TpuCompiler 构造函数会把每个 TPU 编译 phase 一次性注册到全局 phase registry 中。在此二进制中，该钩子是 xla::TpuCompiler::RegisterAllPhases() @ 0xf849ec0。它不是 TpuCompiler_* C-ABI 的一部分，而是编译进 libtpu 的 C++ 方法，因为 XLA 在这里被静态链接；但构建宿主侧的重新实现者必须在任何 compile 前准确调用一次它，这就是为什么它是事实上的 “Initialize”。

注意事项

RegisterAllPhases 也是通向 PJRT PhaseCompile 扩展的桥：RegisterAllPhases 填充的同一 phase 集合，正是 PJRT PhaseCompile 扩展枚举并逐 phase 驱动的集合。该扩展对这项注册的包装由 PJRT PhaseCompile Extension负责；phase 运行的具体 pass 调度由 HLO Pass Registry负责。本页只记录注册入口点存在，并且是 TpuCompiler_New 缺失的 “Initialize” 半边。

函数	地址	角色
`xla::TpuCompiler::RegisterAllPhases`	`0xf849ec0`	一次性填充 phase registry（宿主侧 ctor 钩子）

6. `TpuCompile_*` 支持集群

目的

一个独立的六函数集群，支持 TensorFlow TPUCompileOp kernel，而不是 SE 编译器类。其核心 CompileAndBuild 运行一个完整 TF compile op：输入 metadata proto，驱动 tensorflow::tpu::TpuCompileOpKernelCommon 完成，输出 XLA_TpuProgram 句柄；另外五个是用于 TF 编译缓存和中止策略的小工具。这些函数位于 tpu_util_c_api.cc / 编译缓存键 TU 中，而不是 tpu_executor_c_api.cc。

算法 — 独立 compile

// TpuCompile_CompileAndBuild                                   0xe8bc1e0
// a1=TPUCompileMetadataProto (serialized), a2=mlir/computation input,
// a4/a5 = out program list, a6=StatusRep**
int TpuCompile_CompileAndBuild(meta, input, ..., out_programs, status):
    metadata = parse TPUCompileMetadataProto(meta)
    kernel   = TpuCompileOpKernelCommon(metadata, ...)            // the TF compile-op core
    platform = GetRegisteredDeepseaPlatform()                     // deepsea::executor singleton
    topology = platform->GetTopology()
    result   = kernel.Compile(... topology ...)                   // -> CompiledProgramsAndMetadatas
    if ok:
        for each compiled program in result:
            out_programs[i] = (XLA_TpuProgram*) program           // hand back program handles
    else:
        *status = result.status()

CompileAndBuild 是唯一会产生 TpuProgram 的 TpuCompile_* 函数：它把 TF op 层桥接到序列化 XLA_TpuProgram 句柄，重新实现者预期的 *ToTpuProgram 行为实际上就在这里。

缓存 / 指纹 / 策略辅助函数

// TpuCompile_CreateGuaranteedConstFingerprint                  0xf6a2040
uint64 CreateGuaranteedConstFingerprint(uint64 seed, const char* data, int64 len):
    if len < 0: BUG()                                            // size sanity → trap
    return FingerprintCat2011(seed, Fingerprint2011(data, len))  // 64-bit non-crypto fp

// TpuCompile_DestroyCompilationCacheKey                        0xf6a2e60
void DestroyCompilationCacheKey(void* key, void* prefix):
    if key:    free(key)                                         // two heap strings owned by the key
    if prefix: free(prefix)

// TpuCompile_IsTpuCompilationEnabled                           0xf6a1b40
bool IsTpuCompilationEnabled(): return true                      // constant in this build

// TpuCompile_ShouldTpuCompileOpIgnoreCancellation              0xf6a1b60
bool ShouldTpuCompileOpIgnoreCancellation():
    if !TpuCompilationCancellationTerminatesProcess():
        LOG(WARNING) "...process abort is suppressed... only meant for tests b/79359718"
        return true
    if GetCommandLineOption("xla_jf_exit_process_on_compilation_success") == "true":
        LOG(WARNING) "...abort suppressed when --XLA_jf_exit_process_on_compilation_success... b/72471718"
        return true
    return false

怪异点 — TpuCompile_IsTpuCompilationEnabled 在此构建中是硬编码的 return 1。该函数存在是因为开源 TF op 把它作为运行时 gate 调用，但随附的 libtpu 始终报告 compilation enabled：没有 flag 能关闭它。重新实现者可以在此版本中把该 gate 视作 no-op，但必须保留符号，因为宿主 op 会通过 API 表按名称绑定它。

函数映射

函数	地址	角色	源 TU
`TpuCompile_CompileAndBuild`	`0xe8bc1e0`	运行整个 `TpuCompileOpKernelCommon`；发出 `XLA_TpuProgram` 句柄	`tpu_util_c_api.cc`（HIGH）
`TpuCompile_CreateCompilationCacheKey`	`0xf6a2080`	构建 TF 编译缓存键（config + fingerprints）	cache-key TU
`TpuCompile_CreateGuaranteedConstFingerprint`	`0xf6a2040`	`FingerprintCat2011(seed, Fingerprint2011(data,len))`	cache-key TU
`TpuCompile_DestroyCompilationCacheKey`	`0xf6a2e60`	`free` 缓存键内部的两个堆字符串	cache-key TU
`TpuCompile_IsTpuCompilationEnabled`	`0xf6a1b40`	常量 `true` 运行时 gate	`tpu_util_c_api.cc`
`TpuCompile_ShouldTpuCompileOpIgnoreCancellation`	`0xf6a1b60`	已取消 compile 的中止抑制策略（仅测试）	`tpu_util_c_api.cc`

注意事项

缓存键集群是 XLA TPU 编译缓存的数据路径：CreateCompilationCacheKey 从 metadata 和 guaranteed-const 指纹组装键字符串，CreateGuaranteedConstFingerprint 对 const 输入产生 64 位指纹，DestroyCompilationCacheKey 释放该键拥有的两个堆缓冲区（键本身和 prefix）。Fingerprint2011 是非加密 64 位哈希；重新实现者必须逐位精确复现 Fingerprint2011 和 FingerprintCat2011 combiner，否则缓存键不会在 host/plugin 分割两侧发生相同碰撞。

陷阱 — ShouldTpuCompileOpIgnoreCancellation 中两个中止抑制分支明确是仅测试用（引用 bug refs b/79359718、b/72471718），并且都会在返回 true 前记录一条 warning。在生产中两个谓词均为 false，因此函数返回 false，已取消的 compile 会中止进程。为了“安全”而默认返回 true 的重新实现，会反转 TPU compile op 预期的 fail-fast 行为。

名称	关系
`xla::jellyfish::DeepseaCompiler`	每个 `TpuCompiler_*` 句柄背后的真实 `xla::Compiler` 子类
`xla::TpuCompiler`（宿主侧 shim）	方法转发到此名册的开源类；ctor 调用 `RegisterAllPhases`
`xla::Compiler::Compile` / `RunHloPasses` / `RunBackend`	C-ABI 分发进入的 vtable 方法（`+24`/`+32`）
`ApiConverter::ToC` / `FromC`	跨接缝编组 `XLA_Shape` / `XLA_HloModuleConfig`（见 overview）
`stream_executor::tpu::DeserializeProto<…>`	解包 compile 入口接收的 `TpuSerializedProto`（ptr+len）blob
`deepsea::executor::DeepseaPlatform`	绑定进 `CompileOptions.device_allocator` 的进程级 platform 单例
`WrapperDeviceMemoryAllocator`	compile 入口围绕 `SE_StreamExecutor*` 包装的 `DeviceMemoryAllocator` 子类（vtable `off_21616EF8`）
`tensorflow::tpu::TpuCompileOpKernelCommon`	`TpuCompile_CompileAndBuild` 驱动的 TF op 核心
`XLA_TpuProgram`	`CompileAndBuild` 发出而 `RunBackend` 不发出的序列化程序句柄

交叉引用

The TfTpu C-API Shim — 此名册依赖的 ABI 接缝、*ApiFn() 访问器模式以及不透明句柄 / ApiConverter 约定
TpuProgram 名册 — TpuCompile_CompileAndBuild 产生的 XLA_TpuProgram 序列化程序句柄
TpuExecutable 名册 — TpuCompiler_RunBackend / Compile 返回的装盒 xla::Executable* 句柄
HLO Pass Registry — RunHloPasses 在 DeepseaCompiler vtable 内部执行的具体 pass 调度
PJRT PhaseCompile Extension — 包装 xla::TpuCompiler::RegisterAllPhases 并逐 phase 驱动编译的 PJRT 扩展

TpuCompiler 名册 ​

摘要 ​

1. 生命周期 — New / Free ​

目的 ​

算法 ​

函数映射 ​

注意事项 ​

2. HLO Passes 与 Backend — RunHloPasses / RunBackend ​

目的 ​

入口点 ​

算法 ​

函数映射 ​

注意事项 ​

3. Compile — 融合入口 {#3-compile--the-fused-entry} ​

目的 ​

算法 ​

函数映射 ​

注意事项 ​

4. Shape 查询 — ShapeSize / DefaultDeviceShapeRepresentation ​

目的 ​

算法 ​

函数映射 ​

注意事项 ​

5. Phase 注册 — “Initialize” 步骤 {#5-phase-registration-the-initialize-step} ​

目的 ​

注意事项 ​

6. TpuCompile_* 支持集群 ​

目的 ​

算法 — 独立 compile ​

缓存 / 指纹 / 策略辅助函数 ​

函数映射 ​

注意事项 ​

相关组件 ​

交叉引用 ​

TpuCompiler 名册

摘要

1. 生命周期 — `New` / `Free`

目的

算法

函数映射

注意事项

2. HLO Passes 与 Backend — `RunHloPasses` / `RunBackend`

目的

入口点

算法

函数映射

注意事项

3. Compile — 融合入口 {#3-compile--the-fused-entry}

目的

算法

函数映射

注意事项

4. Shape 查询 — `ShapeSize` / `DefaultDeviceShapeRepresentation`

目的

算法

函数映射

注意事项

5. Phase 注册 — “Initialize” 步骤 {#5-phase-registration-the-initialize-step}

目的

注意事项

6. `TpuCompile_*` 支持集群

目的

算法 — 独立 compile

缓存 / 指纹 / 策略辅助函数

函数映射

注意事项

相关组件

交叉引用