TpuExecutable 名册

本页所有地址均适用于 libtpu-0.0.40-cp314 wheel 中的 libtpu.so（构建 libtpu_lts_20260413_b_RC00，build-id md5 89edbbe81c5b328a958fe628a9f2207d，781,691,048 字节，ELF x86-64 DYN，未 strip）。.text VMA 等于文件偏移。C-ABI 名称（TpuExecutable_*）由 IDA 从调用目标以及烘入 CHECK/MakeErrorImpl 调用点的 tpu_executor_c_api.cc / tpu_execute_c_api.cc 源路径字符串中恢复；反混淆后的 C++ 目标均逐字引用。其他版本会不同。

摘要

TpuExecutable_* 是 已编译 TPU 程序运行形态 的 C-ABI 名册：这是一次编译被反序列化到设备后 StreamExecutor 主机所持有的句柄，主机通过它入队启动、查询程序的 HLO、为缓存生成 fingerprint，并重新序列化程序。TpuProgram_* 负责可序列化、可查询的编译产物，TpuCompiler_* 负责生成它，而 TpuExecutable_* 是覆盖实时 xla::Executable 子树的薄 extern "C" 外壳：主机从不按值看到 xla::TpuExecutable，只看到一个不透明句柄和一组固定的自由函数，这些函数解组 C 结构体并跳转到 C++ executable。

该名册正好有 9 个自由函数，是具名设备运行时函数组中最小的一个。它们清晰分成四个区域：执行（LoadProgramAndEnqueueToStream、ExecuteAsyncOnStream，驱动设备的两个大型编组函数）、序列化（Serialize、Deserialize，通过 xla::DeepseaExecutableProto 往返 proto）、元数据（Fingerprint、HloModule，只读访问器）以及生命周期 / 数组释放（Free、FreeXlaShapeIndexArray、FreeMaybeOwningDeviceAddressArray）。每个入口都是转发器：C-ABI 函数持有一个 XLA_TpuExecutable* 风格的不透明句柄，其第一个槽位指向 C++ 对象，并分派到 xla::TpuExecutable / xla::legacy::TpuExecutableInterface 的方法；具体叶子是 xla::jellyfish::DeepseaExecutable。

C++ 侧是三层类塔。xla::Executable 是上游基类；xla::legacy::TpuExecutableInterface（ExecuteAsyncOnStream @ 0x1342cd20）统一覆盖参数/输出编组；具体的 xla::jellyfish::DeepseaExecutable 实现设备入队（LoadProgramAndEnqueueToStream @ 0x13426260、ToProto @ 0x134282e0、FromProto @ 0x134283e0、fingerprint @ 0x13428a80）。包装类 xla::TpuExecutable（方法集中在 0xf8a7560–0xf8adce0）是 C-ABI 名册逐槽镜像的 legacy StreamExecutor 门面。本页负责逐函数名册和 impl-symbol 映射；不重新推导运行时执行算法。

对重实现而言，契约是：

9 函数名册：名称、impl 符号 + 地址、它跳转到的 C++ 方法，以及它属于四个区域中的哪一个。
句柄约定：每个入口接受一个不透明句柄，其 *handle（或 **handle）是 C++ xla::Executable 子树指针；C-ABI 函数只解组 C 结构体并分派。
序列化线格式：xla::DeepseaExecutableProto 是 wire shape；Serialize = ToProto + proto 拷贝，Deserialize = ParseFromArray + FromProto，二者都返回 StatusOr 风格的 C-ABI 状态对象。
所有权规则：Free 通过槽位 +8 运行 C++ 虚析构函数，然后 free() 句柄盒；两个 FreeArray helper 是裸 free()；HloModule 是唯一会分配的元数据访问器（通过 ApiConverter::ToC 分配一个 XLA_HloModule）。


名册大小	9 个 `extern "C"` `TpuExecutable_*` 函数（已对照函数表验证）
源文件	`tpu_executor_c_api.cc`（大多数）、`tpu_execute_c_api.cc`（`LoadProgramAndEnqueueToStream`）
包装类	`xla::TpuExecutable`（方法 `0xf8a7560`–`0xf8adce0`）— legacy SE 门面
编组基类	`xla::legacy::TpuExecutableInterface::ExecuteAsyncOnStream` @ `0x1342cd20`
具体叶子	`xla::jellyfish::DeepseaExecutable`（`LoadProgram` @ `0x13426260`、`ToProto` @ `0x134282e0`、`FromProto` @ `0x134283e0`、`fingerprint` @ `0x13428a80`）
序列化 proto	`xla::DeepseaExecutableProto`（`operator new(0x28)` = 40 字节）
句柄编组	`HloModule` 使用 `ApiConverter::ToC(xla::HloModule const&)`
证据等级	可重实现级 / 经 IDA 反编译逐字节确认

范围 — 框定本名册的不透明句柄 + *ApiFn() 访问器模型（主机如何通过函数指针槽位而不是符号触达 TpuExecutable_Fingerprint）见 TfTpu C-API Shim 概览。此 executable 包装的编译程序句柄见 TpuProgram 名册。两个大型入口背后的运行时执行算法归属于流上的异步执行和加载程序并入队；本页是这些入口之上的 C-ABI 名册，不是它们的第二份算法副本。与这条 legacy 路径并存且不通过它路由的 PJRT executable 生命周期见 PJRT Executable 执行。

1. 名册概览 {#1-the-roster-at-a-glance}

九个函数按区域分组，列出 impl 符号（始终是 C-ABI 自由函数自身）、地址和大小，以及它分派到的 C++ 方法。置信度反映反编译函数体中跳转目标被观察到的直接程度。

`TpuExecutable_*`	地址	大小	跳转到（C++）	区域
`LoadProgramAndEnqueueToStream`	`0xeaafba0`	4496	`xla::jellyfish::DeepseaExecutable::LoadProgramAndEnqueueToStream` @ `0x13426260`（通过 vtable+96）	执行
`ExecuteAsyncOnStream`	`0xeabd500`	4708	`xla::legacy::TpuExecutableInterface::ExecuteAsyncOnStream` @ `0x1342cd20`（通过 vtable+24）	执行
`Serialize`	`0xeabea80`	178	`xla::jellyfish::DeepseaExecutable::ToProto` @ `0x134282e0`	序列化
`Deserialize`	`0xeabede0`	288	`xla::jellyfish::DeepseaExecutable::FromProto` @ `0x134283e0`	序列化
`Fingerprint`	`0xeabea40`	54	`DeepseaExecutable::fingerprint` @ `0x13428a80`（缓存字段位于 `obj+96` → `+648/+656`）	元数据
`HloModule`	`0xeabef00`	86	对 `executable.hlo_module_` 执行 `ApiConverter::ToC(xla::HloModule const&)`	元数据
`Free`	`0xeabef60`	51	通过 vtable `+8` 调用 C++ 虚析构函数，然后 `free()`	生命周期
`FreeXlaShapeIndexArray`	`0xeabea00`	10	裸 `free(ptr)`	生命周期
`FreeMaybeOwningDeviceAddressArray`	`0xeabea20`	10	裸 `free(ptr)`	生命周期

NOTE — PJRT_TpuExecutable_* 族（PJRT_TpuExecutable_RunHloCostAnalysis_Args、..._GetCompiledMemoryStats_Args、..._GetHloModuleWithConfig_Args、..._SetTpuCompilationEnv_Args 等）是不同表面，不在本名册中。这些是由 pjrt::(anonymous namespace)::RunHloCostAnalysis / GetCompiledMemoryStats / 等处理的 PJRT-extension 参数结构体，属于 PJRT 页面。本文负责的 C-ABI 名册正好是动态符号名以 TpuExecutable_ 开头（没有 PJRT_ 前缀）的九个函数。按前缀过滤是唯一可靠的分隔方式：两族共享动词（HloModule、cost analysis），但位于 ABI 的相反两侧。
QUIRK — 包装类名为 xla::TpuExecutable（因此 TpuExecutable_HloModule ↔ 读取 hlo_module_ 的成员），但执行入口会比 xla::TpuExecutable 自身方法更深一到两层：ExecuteAsyncOnStream 落到 xla::legacy::TpuExecutableInterface 基类（实现编组的类），LoadProgramAndEnqueueToStream 落到 xla::jellyfish::DeepseaExecutable 叶子。若重实现者把每个 TpuExecutable_* C 函数 1:1 映射为一个 xla::TpuExecutable 方法，就会放错两个大型入口；它们的目标是 interface/leaf，而不是 wrapper。

2. 执行入口

目的

两个执行函数是本名册中唯一的大函数体（4496 B 和 4708 B）；其他七个合计不到 700 B。二者都是 C-ABI 编组层：从 SE shim 接收扁平 launch/run 结构，解组为设备运行时期望的类型化 C++ 对象（xla::ExecutableRunOptions、xla::ExecutionInput 向量、DeviceAddressBase span、一个反序列化的 DeviceAssignment），通过 vtable 槽位分派到实际入队逻辑，并把结果（ExecutionOutput / absl::Status）重新编组为 C 结构体交还调用方。它们都不包含设备入队算法；它们只是算法之上的 C 边界。

入口点

text

TpuExecutable_LoadProgramAndEnqueueToStream   0xeaafba0   (tpu_execute_c_api.cc)
  └─ un-marshal C structs → dispatch through the handle vtable
       └─ vtable+96 ──► xla::jellyfish::DeepseaExecutable::LoadProgramAndEnqueueToStream
                          0x13426260  (deepsea_executable.cc) ── real load+enqueue

TpuExecutable_ExecuteAsyncOnStream            0xeabd500   (tpu_executor_c_api.cc)
  └─ vtable+24 ──► xla::legacy::TpuExecutableInterface::ExecuteAsyncOnStream
                     0x1342cd20  ── arg/output marshal + dispatch
                       └─ (leaf) DeepseaExecutable::LoadProgramAndEnqueueToStream  0x13426260
```text

### 算法

C-ABI 外壳是统一的；设备工作位于运行时页面。`ExecuteAsyncOnStream` 的外壳形态如下：

```c
// TpuExecutable_ExecuteAsyncOnStream                          0xeabd500
function TpuExecutable_ExecuteAsyncOnStream(handle /*a1*/, run_opts /*a2*/,
                                            inputs /*a3*/, n_inputs /*a4*/,
                                            out /*a5*/, status /*a6*/):
    exe   = *handle                          // opaque box → xla::Executable subtree
    // un-marshal C structs into ExecutableRunOptions + vector<ExecutionInput>
    cpp_opts   = FromC(run_opts)
    cpp_inputs = FromC(inputs, n_inputs)
    // dispatch through the executable vtable slot +24 (ExecuteAsyncOnStream)
    result = (*exe->vtable[+24])(exe, &cpp_opts, move(cpp_inputs))   // 0x1342cd20
    // re-marshal ExecutionOutput / Status back to C
    if result.ok: ToC(result.value, out)
    else:         set_status(status, result.status)

LoadProgramAndEnqueueToStream 遵循同一模式，但通过 vtable 槽位 +96 分派到 DeepseaExecutable 叶子，并且解组步骤更重：它重建一个 HloModule（仅用于 shape signature）、一个 ExecutableRunOptions、一个反序列化的 DeviceAssignment，以及两个 DeviceAddressBase 向量。这两个函数体的完整逐行重建，包括 absl::Status 处理和设备侧 DeepseaExecutor::LoadProgram → DeepseaStream::EnqueueRequest 链，位于运行时页面。

函数映射

函数	大小	角色
`TpuExecutable_LoadProgramAndEnqueueToStream` @ `0xeaafba0`	4496	C-ABI 编组 → `DeepseaExecutable::LoadProgramAndEnqueueToStream`（vtable+96）
`TpuExecutable_ExecuteAsyncOnStream` @ `0xeabd500`	4708	C-ABI 编组 → `TpuExecutableInterface::ExecuteAsyncOnStream`（vtable+24）

注意事项

这两个入口是进入同一个设备运行时的两扇公共门。ExecuteAsyncOnStream 由 xla::LocalClient / xla::Service StreamExecutor 栈触达（LocalExecutable::RunAsync → Executable::ExecuteAsyncOnStreamWrapper → vtable+24）；C-ABI shim 是第二扇门，它解组 C 结构体并调用同一个 +24 槽位。LoadProgramAndEnqueueToStream 是更低层的入队原语，interface 的 ExecuteAsyncOnStream 自身最终也会通过叶子驱动它。现代 PJRT 路径（基于 tpu::System 的 xla::TpuClient）不经过二者；它通过 tpu::System::Execute（0x1d0b33e0）触达同一个设备运行时。重实现 legacy SE 后端时应复现这两个 C 函数；重实现 PJRT 时则忽略它们。

3. 序列化入口

目的

Serialize 和 Deserialize 通过 xla::DeepseaExecutableProto（wire form）往返 executable。Serialize 生成一个堆拥有的 proto，调用方稍后释放；Deserialize 把字节缓冲区解析成新的 executable，并返回一个 C-ABI 状态对象。二者都很短（178 B / 288 B），因为真实工作在 DeepseaExecutable::ToProto / FromProto C++ 方法中；C 函数只是 proto 分配 + 状态管道。

算法

// TpuExecutable_Serialize                                     0xeabea80
function TpuExecutable_Serialize(self /*_XMM0/a1*/, out_proto /*a3*/):
    *out_proto = NULL
    proto = operator new(0x28)                  // 40-byte DeepseaExecutableProto
    DeepseaExecutableProto::DeepseaExecutableProto(proto, 0)
    *out_proto = proto
    tmp = DeepseaExecutable::ToProto(self)       // 0x134282e0 — fill a stack proto
    // move-or-copy tmp into the heap proto, fast-pathing the arena-equal case
    if tmp.arena == proto.arena: InternalSwap(proto, tmp)   // same arena → swap
    else:                        CopyFrom(proto, tmp)       // cross-arena → deep copy
    ~DeepseaExecutableProto(tmp)
```text

```c
// TpuExecutable_Deserialize                                   0xeabede0
function TpuExecutable_Deserialize(len /*a1*/, bytes /*a2*/, out /*a3*/, status /*a4*/):
    proto = DeepseaExecutableProto(0)                          // stack proto
    if !MessageLite::ParseFromArray(proto, bytes, len):
        set_status(status,
            MakeErrorImpl<INTERNAL>(                            // code 13
                "TpuExecutable_Deserialize: proto deserialization failed"))
        goto cleanup
    exe = DeepseaExecutable::FromProto(proto)                  // 0x134283e0 → StatusOr
    if exe.ok():
        box = operator new(8)                                  // 8-byte handle box
        *box = exe.value;  *out = box                          // hand off ownership
    else:
        set_status(status, exe.status)
        if exe.value: (*exe.value->vtable[+8])(exe.value)      // destroy partial leaf
cleanup:
    ~DeepseaExecutableProto(proto)

GOTCHA — 反序列化后的句柄是一个 8 字节堆盒（operator new(8u)），其唯一槽位保存 C++ DeepseaExecutable*。这正是 TpuExecutable_Free 所期望的盒形态：*handle 是 C++ 对象，盒自身由 free() 释放。若重实现返回裸 C++ 指针而不是把它装盒，Free 会把非盒解引用并破坏堆。这个盒故意多一层指针间接：它让主机持有不透明 XLA_TpuExecutable*，同时插件把真实对象放在后面。
NOTE — Deserialize 通过 a4 返回标准 SE-shim 状态对象（一个 refcounted absl::status_internal::StatusRep**），proto 解析失败时 code 13 = INTERNAL。反编译函数体主要由 StatusRep::Unref / _InterlockedIncrement 引用计数搬运组成，而不是反序列化逻辑；后者完全位于 DeepseaExecutable::FromProto 内部。

函数映射

函数	大小	角色
`TpuExecutable_Serialize` @ `0xeabea80`	178	`ToProto` + arena-aware swap/copy 到堆上的 `DeepseaExecutableProto`
`TpuExecutable_Deserialize` @ `0xeabede0`	288	`ParseFromArray` + `FromProto`；把结果装盒，设置 `StatusRep` out-param

注意事项

proto 类型 xla::DeepseaExecutableProto 是 executable 序列化，区别于支撑 TpuProgram_* 句柄的 tensorflow::TPUExecutableProto（后者是馈入编译缓存的 program 序列化）。这里的 Serialize 为已加载 executable 的传输/缓存生成 executable proto；Deserialize 从字节重建它。Serialize 中的 arena 相等快速路径（InternalSwap vs CopyFrom）是 proto2 惯用法：当来源和目的共享 arena 时，移动是指针交换，否则是深拷贝。两条分支都要复现：跨不同 arena swap 会别名到已释放内存。

4. 元数据入口

目的

两个只读访问器。Fingerprint 返回 executable 的缓存 fingerprint（一个 string/Cord 对）以用于 cache key；HloModule 把 executable 的 xla::HloModule 编组到 XLA_HloModule C 结构体中，供主机检查。Fingerprint 不分配；它返回指向 executable 自身字段的指针。HloModule 是唯一会分配的元数据函数（通过 ApiConverter::ToC）。

算法

// TpuExecutable_Fingerprint                                   0xeabea40
function TpuExecutable_Fingerprint(handle /*a1*/, out_ptr /*a2*/, out_len /*a3*/):
    fp = *(*handle + 96)            // DeepseaExecutable cached fingerprint field
    if fp.is_long_form(fp+671 >= 0):      // SSO discriminator byte
        *out_len = fp+671                 // inline (short) string: len byte, data at fp+648
        *out_ptr = fp+648
    else:
        *out_len = *(fp+656)              // heap string: length at +656
        *out_ptr = *(fp+648)              //              data ptr at +648
    return *out_ptr
```text

```c
// TpuExecutable_HloModule                                     0xeabef00
function TpuExecutable_HloModule(out /*a1*/, handle /*a2*/):
    mod = *(*handle + 8)                  // xla::Executable::hlo_module_
    if mod == NULL:
        LOG(FATAL) "hlo_module_ != nullptr"  // executable.h:335
    ApiConverter::ToC(out, mod)           // fill caller-provided XLA_HloModule
    return out

QUIRK — Fingerprint 直接从 C++ 对象偏移 +96 处读取缓存 fingerprint，并内联解码 std::string 的 short-string-optimization 布局（+671 字节的符号区分 inline 与 heap，长度位于 +671 或 +656，数据位于 +648 或 *(+648)）。它不会在调用时调用 DeepseaExecutable::fingerprint（0x13428a80）；该方法填充字段，而此访问器只暴露字段。每次 Fingerprint 调用都重新计算 fingerprint 的重实现也能正确，但会更慢；二进制读取预计算字段并返回借用指针（因此调用方不得释放它们）。
GOTCHA — 如果 hlo_module_（位于 executable+8）为空，HloModule 会 LOG(FATAL)，而不是返回错误状态。只有 executable 在启用 module retention 的情况下构建时，HLO module 才会被保留；没有保留它的 executable 在这里会崩溃，而不是优雅失败。fatal-log 源锚点是 executable.h:335。这是上游 XLA 行为原样通过 C ABI 暴露出来。

函数映射

函数	大小	角色
`TpuExecutable_Fingerprint` @ `0xeabea40`	54	借用 `obj+96` 处缓存 fingerprint 字段的 `(ptr,len)`；SSO 解码
`TpuExecutable_HloModule` @ `0xeabef00`	86	把 `hlo_module_` 通过 `ApiConverter::ToC` 写入调用方 `XLA_HloModule`；为空则 FATAL

注意事项

Fingerprint 返回指向 executable 的借用指针；没有配对 free，数据只在句柄存活期间有效。HloModule 是唯一的分配者：ApiConverter::ToC(xla::HloModule const&) 写入调用方提供的 XLA_HloModule out-struct（作为前导 sret 参数传入），并填充堆拥有的内部内容（见 shim 概览中的编组约定），因此主机必须为 XLA_HloModule 运行匹配的 Destroy/free。相关的更丰富元数据访问器（GetCompiledMemoryStats、GetCostAnalysis、GetOutputLayouts 等）是 xla::TpuExecutable（0xf8a7560–0xf8aab40）上的 C++ 方法，通过 PJRT extension 表面触达，而不是通过本 C-ABI 名册；见 §1 中的说明。

5. 生命周期和数组释放入口

目的

三个释放器。Free 拆毁完整 executable 句柄（析构 C++ 对象，然后释放盒）；两个 FreeArray helper 释放本名册和兄弟页面上的其他函数作为裸堆指针交回调用方的数组。

算法

// TpuExecutable_Free                                          0xeabef60
function TpuExecutable_Free(box /*ptr*/):
    if box != NULL:
        if *box != NULL:
            (*(*box)->vtable[+8])(*box)      // C++ virtual destructor (delete-style)
        free(box)                            // free the 8-byte handle box

// TpuExecutable_FreeXlaShapeIndexArray                        0xeabea00
function TpuExecutable_FreeXlaShapeIndexArray(p):  if p: free(p)

// TpuExecutable_FreeMaybeOwningDeviceAddressArray             0xeabea20
function TpuExecutable_FreeMaybeOwningDeviceAddressArray(p):  if p: free(p)
```text

> **NOTE —** 两个 array-free helper 逐字节相同（`if (a1) free(a1);`，分别是相邻地址 `0xeabea00` / `0xeabea20` 处的 10 字节）。它们故意保留为两个不同符号：每个都与 ABI 中其他位置的不同数组生产调用配对（一个 `XLA_ShapeIndex[]` 和一个 `SE_MaybeOwningDeviceAddress[]`），因此重实现者即使实现同样是 `free`，也必须保留两个名称。把它们折叠成一个函数会破坏 ABI 的具名配对契约，而不是运行时行为。
>
> **GOTCHA —** `TpuExecutable_Free` 期望*装盒*句柄：`box` 是 `Deserialize` 生成的 `operator new(8)` 盒，`*box` 是 C++ `DeepseaExecutable*`，虚析构函数通过 `(*box)->vtable[+8]` 调用。顺序固定：先经 vtable 槽位析构对象，*再* `free()` 盒。先 `free(box)` 会泄漏 C++ 对象及其所有设备侧状态（已加载程序、输出缓冲区）；对盒而不是 `*box` 调用析构函数则会在垃圾上调用析构。

### 函数映射

| 函数 | 大小 | 角色 |
|---|---|---|
| `TpuExecutable_Free` @ `0xeabef60` | 51 | 通过 vtable+8 虚析构 `*box`，然后 `free(box)` |
| `TpuExecutable_FreeXlaShapeIndexArray` @ `0xeabea00` | 10 | `if (p) free(p)` — 释放 `XLA_ShapeIndex[]` |
| `TpuExecutable_FreeMaybeOwningDeviceAddressArray` @ `0xeabea20` | 10 | `if (p) free(p)` — 释放 `SE_MaybeOwningDeviceAddress[]` |

### 注意事项

`Free` 是整个名册的析构器：每个由 `Deserialize`（以及由 compiler/program 路径生成已加载 executable 的代码）铸造的句柄都通过它回收。两个 array free 存在是因为 C ABI 无法返回拥有所有权的 C++ 容器；概念上返回 `vector<ShapeIndex>` 的函数会返回一个裸 `XLA_ShapeIndex*` + count 和一个具名 free，这与 [`TpuProgram_*`](tpu-program-roster.md) 和 [`TpuTopology_*`](tpu-topology.md) 上的 `*FreeArray` 模式一致。每个 producer 都应与自己的具名 free 配对；不要交叉使用（用 device-address helper 释放 `ShapeIndex[]` 仍然只是 `free`，但具名契约记录意图，并且可能在其他构建中分化）。

---

## 相关组件

| 名称 | 关系 |
|---|---|
| `xla::TpuExecutable`（`0xf8a7560`–`0xf8adce0`） | 本名册所镜像的 legacy StreamExecutor 包装类；也是仅 PJRT 元数据方法的归属处 |
| `xla::legacy::TpuExecutableInterface::ExecuteAsyncOnStream` @ `0x1342cd20` | `TpuExecutable_ExecuteAsyncOnStream` 分派进入的编组基类（vtable+24） |
| `xla::jellyfish::DeepseaExecutable`（leaf） | 具体实现：`LoadProgram` `0x13426260`、`ToProto` `0x134282e0`、`FromProto` `0x134283e0`、`fingerprint` `0x13428a80` |
| `xla::DeepseaExecutableProto`（40 字节） | `Serialize` / `Deserialize` 背后的序列化 wire form |
| `ApiConverter::ToC(xla::HloModule const&)` | 为 `TpuExecutable_HloModule` 把 `hlo_module_` 编组进 `XLA_HloModule` C 结构体 |
| `TpuProgram_*` roster | 此*运行中* executable 所包装的可序列化*编译程序*句柄 |

## 交叉引用

- [TfTpu C-API Shim 概览](overview.md) — 框定本名册的不透明句柄 + `*ApiFn()` 访问器模型（三跳 accessor → slot → C-ABI impl 模式）
- [TpuProgram 名册](tpu-program-roster.md) — 此 executable 所包装的 `TpuProgram_*` 编译程序句柄（基于 `tensorflow::TPUExecutableProto`）
- [TpuCompiler 名册](tpu-compiler-roster.md) — 生成加载到这些 executable 中的 program 的 `TpuCompiler_*` / `TpuCompile_*` 表面
- [TpuExecutor 名册](tpu-executor-roster.md) — executable 入队到的逐设备 `TpuExecutor_*` 运行时
- [流上的异步执行](../runtime/execute-async-on-stream.md) — `TpuExecutable_ExecuteAsyncOnStream` 背后的运行时算法（本页是其上的 C-ABI 名册）
- [加载程序并入队](../runtime/load-program-enqueue.md) — `TpuExecutable_LoadProgramAndEnqueueToStream` 背后的运行时算法（load + device enqueue）
- [PJRT Executable 执行](../pjrt/executable-execution.md) — 与此 legacy SE 路径并存且不通过它路由的现代 PJRT executable 生命周期

TpuExecutable 名册 ​

摘要 ​

1. 名册概览 {#1-the-roster-at-a-glance} ​

2. 执行入口 ​

目的 ​

入口点 ​

函数映射 ​

注意事项 ​

3. 序列化入口 ​

目的 ​

算法 ​

函数映射 ​

注意事项 ​

4. 元数据入口 ​

目的 ​

算法 ​

函数映射 ​

注意事项 ​

5. 生命周期和数组释放入口 ​

目的 ​

算法 ​

TpuExecutable 名册

摘要

1. 名册概览 {#1-the-roster-at-a-glance}

2. 执行入口

目的

入口点

函数映射

注意事项

3. 序列化入口

目的

算法

函数映射

注意事项

4. 元数据入口

目的

算法

函数映射

注意事项

5. 生命周期和数组释放入口

目的

算法