TfTpu C-API Shim

本页所有地址均适用于来自 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；反修饰后的 C++ 符号和导出的 C-ABI 符号名均按原样引用）。.text VMA 等于文件偏移。其他版本会有所不同。

摘要

TfTpu C-API shim 是一个稳定的 C ABI，开源 StreamExecutor TPU 后端，也就是随 TensorFlow/XLA 内部发布、而非位于此二进制内部的 C++ 代码，通过它驱动 libtpu.so，同时完全不链接任何 libtpu C++ 符号。它存在的原因只有一个：闭源插件和开源宿主不能跨 .so 边界共享 std::string、absl::Status、xla::Shape 或任何其他 C++ 类型，因为两者由不同工具链构建，且没有 ABI 契约。在二者之间铺设一个扁平 C 接口后，耦合就消失了。这个 shim 就是该接口：数百个 extern "C" 自由函数，按其支撑的 C++ 类分组，只接收和返回不透明结构句柄与 POD。

该表面有三个结构性部分，本页负责为三者提供导览。第一，按类划分的 C 函数名册：TpuCompiler_*、TpuExecutable_*、TpuExecutor_*、TpuTransferManager_*、TpuProgram_*、TpuPlatform_*、TpuTopology_*、TpuEmbeddingEngine_*、TpuConfigurationApi_*（以及较小的 TpuCoreLocation_*、TpuNodeContext_*、TpuMeshState_*、TpuCompile_* 簇）。每个簇都是一个 stream-executor 抽象的 C 侧面，并有专门的同级页面清点其函数。第二，*ApiFn() 访问器模式：宿主不会按动态符号名调用 TpuExecutor_Allocate。它通过一个单例的函数指针结构体调用，stream_executor::tpu::ExecutorApiFn() 返回指向函数内静态 TfTpu_ExecutorApiFn 的指针，该结构体成员在初始化时填充一次，每个 SE shim 方法索引该结构体的一个槽位。第三，不透明句柄约定：每个跨边界的 C++ 对象都由一个 SE_* / XLA_* / Tpu* 不透明 C 结构体镜像，ApiConverter::ToC / FromC / Destroy 在丰富的 C++ 类型与其扁平 C 双胞胎之间编组。

这是一个地图页。它建立 ABI 模型和访问器模式，然后用一两句话概述每个名册，并链接到包含逐函数细节的同级页面。它不重复逐函数表格、表填充引导过程（归属于 TfTpu_Initialize 引导），也不重复消费 ExecutorApiFn 表的 SE platform/executor 对象模型（归属于 StreamExecutor Platform 与 Executor 模型）。这是重新实现者在深入任何单个名册之前首先阅读的导览。

对于重新实现，契约是：

ABI 边界本身：为什么存在扁平 C 接口，什么会跨过它（只有不透明句柄 + POD），以及什么绝不会跨过它（没有按值传递的 std::/absl::/xla:: 类型）。
*ApiFn() 访问器模式：三个访问器（ExecutorApiFn、OpsApiFn、ProfilerApiFn），各自返回由初始化引导填充的单例函数指针结构体；shim 方法如何解引用槽位；IsStreamExecutorEnabled 如何在使用前探测表。
不透明句柄约定：SE_* / XLA_* / Tpu* 结构体族，以及 ApiConverter::ToC / FromC / Destroy 编组，并列出具体的 C++↔C 类型配对。
名册地图：哪个 C 函数簇支撑哪个 SE 抽象，以及各自记录在何处。


ABI 边界	`extern "C"` 自由函数；只有不透明句柄 + POD 跨越 `.so` 边界
访问器（executor/stream）	`stream_executor::tpu::ExecutorApiFn() @ 0x20819360` → `&ExecutorApiFn()::executor_api_fn`（函数内静态）
访问器（embedding/ops）	`stream_executor::tpu::OpsApiFn() @ 0x10900e80` → `&OpsApiFn()::ops_api_fn`
访问器（profiler）	`stream_executor::tpu::ProfilerApiFn() @ 0x10900ea0` → `&ProfilerApiFn()::profiler_api_fn`
启用探测	`stream_executor::tpu::IsStreamExecutorEnabled(TfTpu_ExecutorApiFn*) @ 0x20819380`
初始化探测	`stream_executor::tpu::IsInitialized(TfTpu_ExecutorApiFn*) @ 0x208193c0`
不透明句柄编组	`ApiConverter::ToC`（9 个重载）/ `FromC`（6 个）/ `Destroy`（3 个）
C-ABI 名册（此二进制侧）	9 个命名簇，约 150 个自由函数（`TpuCompiler_` 7 个，`TpuExecutor_` 25 个，`TpuTransferManager_` 19 个，`TpuProgram_` 18 个，`TpuTopology_` 17 个，`TpuEmbeddingEngine_` 15 个，`TpuPlatform_` 11 个，`TpuExecutable_` 9 个，`TpuConfigurationApi_*` 8 个）
共存的现代 ABI	PJRT C-API（`PJRT_*`）——见范围说明
证据等级	重新实现级 / 由 IDA 反编译按字节确认

范围 — *ApiFn 结构体的一次性填充（哪个符号解析到哪个槽位，以及何时发生）归属于 TfTpu_Initialize 引导。ExecutorApiFn 的消费者，也就是通过该表转发每个方法的 TpuPlatform / TpuExecutor SE 对象模型，归属于 StreamExecutor Platform 与 Executor 模型。与这个遗留 SE shim 共存的现代 PJRT C-API 归属于 PJRT 概览。本页只覆盖 shim-ABI 导览 + 访问器模式 + 句柄约定 + 名册地图；它会链接到那些页面，而不是重述它们。

1. 为什么存在 C Shim

ABI 边界

libtpu.so 是一个闭源 Google 构建；StreamExecutor TPU 后端（stream_executor::tpu::TpuExecutor、tensorflow::tpu::TpuPlatform、xla::TpuCompiler 及相关组件）是开源的，由编译周围 TensorFlow/XLA 的任意工具链构建。如果这两半按值传递 C++ 对象，例如 xla::Shape、absl::Status、std::vector<...>，它们就必须逐位一致地约定每个标准库和 XLA 类型的布局。它们没有也不可能做到：插件在构建时冻结，而宿主独立演进。扁平 C 接口消除了这个问题。C 拥有稳定调用约定且没有名称修饰，因此声明为 extern "C" 的函数可跨边界调用，无论两侧由哪个编译器生成。

具体来说，没有任何带 C++ ABI 的东西跨越边界。 跨越边界的是：

不透明句柄：指向调用方永不检查的结构体的指针（SE_StreamExecutor*、XLA_Shape*、SE_Event*、TpuExecutor*）。宿主持有它们、传回它们，并通过配对的 _Free 调用释放它们。
POD：整数、大小、原始 const char* + 长度配对，以及小型固定 C 结构体。
status-out 临时对象：被调用方填充、调用方查询 ok/code/message 后释放的 C TF_Status 风格对象（SE 页面记录的 scratch / op / ok?-code-msg / free 惯用法）。

两侧并不是通过动态链接器在加载时解析 TpuExecutor_Allocate 连接起来，而是通过插件早期一次性交给宿主的函数指针表连接起来。该表就是 *ApiFn 结构体。

两半分别位于何处

IDA 反编译让这种拆分可见。被调用方半边，即真正触碰 TPU 驱动的实现，作为 extern "C" 自由函数位于 libtpu.so 中，其 IDA 恢复出的名称带有各名册前缀：TpuExecutor_Allocate @ 0xeab9120、TpuCompiler_New @ 0xeabc4a0、TpuProgram_New @ 0xe8bda60、TpuTransferManager_CanBufferBeAccessedNow @ 0xeaba7a0，以及约 150 个其他函数。调用方半边，即把每个 SE 虚方法转发到 *ApiFn 表槽位的薄 C++ shim，也作为经修饰的 stream_executor::tpu::TpuExecutor::* / tensorflow::tpu::TpuTransferManager::* / xla::TpuCompiler::* 方法出现在此二进制中，因为 XLA 被静态链接进了 libtpu.so。重新实现宿主侧时要复现第二半；重新实现插件时要复现第一半。

注意 — C-ABI 自由函数是导出的、带版本标签的动态符号，名称为 TpuCompiler_*、TpuExecutor_* 等（每个都是 @@VERS_1.0，类型 T/DF .text），而不是 TfTpu_Compiler_*；不存在 TfTpu_<Class>_* 符号。TfTpu_ 前缀只保留在表类型名（TfTpu_ExecutorApiFn）和三个公共 init/server 入口点（TfTpu_Initialize、TfTpu_InitializeTpuModelServer、TfTpu_GetTpuPartitionedCallParams）上。应把 Tpu<Class>_<Method> 视为此构建中的规范 C-ABI 名册名称。

2. `*ApiFn()` 访问器模式

单例函数指针结构体

宿主从不直接命名 C-ABI 函数。每个名册都通过从零参数访问器取得的单例函数指针结构体访问。有三个这样的访问器，每个都返回一个函数内静态结构体实例的地址：

// stream_executor::tpu::ExecutorApiFn()                        sub_20819360
void* ExecutorApiFn():
    return &ExecutorApiFn()::executor_api_fn   // function-local static TfTpu_ExecutorApiFn

// stream_executor::tpu::OpsApiFn()                             sub_10900E80
void* OpsApiFn():
    return &OpsApiFn()::ops_api_fn

// stream_executor::tpu::ProfilerApiFn()                        sub_10900EA0
void* ProfilerApiFn():
    return &ProfilerApiFn()::profiler_api_fn
```text

每个访问器都是返回静态结构体稳定指针的 Meyers 单例。结构体成员是带类型的函数指针，每个 C-ABI 名册条目一个，按固定字节偏移布局。SE shim 方法按偏移而非按名称寻址这些成员；SE 页面记录 executor 侧槽位模式（`ExecutorApiFn()[+16]`、`[+32]`、`[+360]` 等），其中每个偏移都是 `TfTpu_ExecutorApiFn` 的一个成员。重新实现者必须让*填充*表的一侧与*读取*表的一侧保持槽位偏移完全相同，因为偏移就是整个契约，调用时没有名称可回退。

三个访问器按子系统而非每类一表来划分表面：

| 访问器 | 地址 | 支撑 | 承载的名册簇 |
|---|---|---|---|
| `ExecutorApiFn` | `0x20819360` | 设备运行时：platform、executor、stream、transfer、program、compiler | `TpuExecutor_*`、`TpuTransferManager_*`、`TpuPlatform_*`、`TpuProgram_*`、`TpuCompiler_*`、`TpuExecutable_*`、`TpuTopology_*`、`TpuCoreLocation_*`、`TpuNodeContext_*`、`TpuConfigurationApi_*`、`TpuMeshState_*`、`TpuCompile_*` |
| `OpsApiFn` | `0x10900e80` | Embedding / SparseCore ops | `TpuEmbeddingEngine_*` 和 XLA-op-kernel embedding 表面 |
| `ProfilerApiFn` | `0x10900ea0` | Profiler / trace collection | TPU profiler C 表面 |

> **特性 —** 按类划分的名册拆分是*文档*结构，而不是*表*结构。二进制把所有设备运行时类（`Compiler`、`Executor`、`Stream`、`Program`、`TransferManager`、`Platform`、`Topology` 等）分组到一个由单个访问器到达的 `TfTpu_ExecutorApiFn` 结构体中；embedding 和 profiler 各自有表。若重新实现者为每个 C++ 类构建一个结构体，就不会匹配二进制的偏移布局。这里有三张表，而不是九张。

### 使用前探测表

由于结构体是函数内静态对象，在引导填充它之前，其成员为零。shim 用两个直接接收表指针的谓词防护未填充表：

```c
// stream_executor::tpu::IsStreamExecutorEnabled(TfTpu_ExecutorApiFn* a1)   sub_20819380
char IsStreamExecutorEnabled(a1):
    if (!a1->slot[0]):              return 0    // first fn-ptr unset → table not populated
    enabled = a1->slot[0](a1)                   // call the "is enabled" entry
    if (!enabled):                  return 0
    a1->slot[8](enabled)                        // consume/dispose the returned probe object
    return 1

IsStreamExecutorEnabled（0x20819380）及其同级 IsInitialized（0x208193c0）都接收 TfTpu_ExecutorApiFn* 并解引用其前导槽位。第一个槽位同时作为空检查是有意的：全零表会被读作“SE 后端禁用”，每个受门控路径（例如 SE 页面显示的、由 IsStreamExecutorEnabled(ExecutorApiFn()) 门控的 RegisterTpuPlatform）都会把它当成干净的 no-op 而非崩溃。重新实现者必须让第一个成员成为有效的谓词函数，并且必须将整个结构体零初始化，使未引导的表能安全地探测失败。

陷阱 — 探测会调用 slot[0]，然后调用 slot[8]，把 slot[0] 返回的值视为由 slot[8] 释放的拥有对象。若重新实现让 slot[0] 返回借用/静态指针，再通过 Free 槽位释放，就会双重释放或释放静态对象；若让 slot[0] 在“enabled”时返回 nullptr，就会报告后端禁用。成功时谓词必须返回一个新的、可释放对象。

访问器返回单例的指针，所以调用方既能读取槽位，也能在引导期间通过同一地址写入槽位。填充槽位的引导过程是 TfTpu_Initialize 引导；本页只确认访问器和探测存在。

3. 不透明句柄约定

`SE_` / `XLA_` / `Tpu*` 镜像

每个必须跨越边界的 C++ 对象都有一个扁平 C 双胞胎。命名是一致的：stream_executor::Foo 变为 SE_Foo，xla::Foo 变为 XLA_Foo，TPU 专用句柄保留其 Tpu* 名称。C 侧把这些当作不透明对象：它持有指针，把指针传给下一次调用，并且从不读取字段。丰富形式与扁平形式之间的编组由 ApiConverter 完成，包含三个动词：

ToC：把 C++ 对象序列化进其 C 结构体（输出参数或按值句柄）。观察到 9 个重载（加上 ToC(DeviceAddressAllocator*) 内部的两个 lambda thunk）。
FromC：从 C 结构体重建 C++ 对象。观察到 6 个重载。
Destroy：释放 C 结构体的堆拥有内部数据（XLA_Shape、XLA_Literal、XLA_ShapedBuffer）。观察到 3 个重载。

从符号表恢复出的具体 C++↔C 配对：

C++ 类型	C 结构体（句柄）	编组动词
`xla::Shape`	`XLA_Shape`	`ToC(const Shape&, XLA_Shape)`、`FromC(const XLA_Shape)`、`Destroy(XLA_Shape*)`
`xla::Layout`	`XLA_Layout`	`ToC(const Layout&, XLA_Layout)`、`FromC(const XLA_Layout)`
`xla::LiteralSlice` / `Literal`	`XLA_Literal`	`ToC(const LiteralSlice&, XLA_Literal)`、`FromC(XLA_Literal)`、`Destroy(XLA_Literal*)`
`xla::ShapedBuffer`	`XLA_ShapedBuffer`	`ToC(const ShapedBuffer&, XLA_ShapedBuffer)`、`FromC(XLA_ShapedBuffer)`、`Destroy(XLA_ShapedBuffer*)`
`xla::HloModule`	（C module 句柄）	`ToC(const HloModule&)`
`xla::HloModuleConfig`	`XLA_HloModuleConfig`	`ToC(const HloModuleConfig&)`、`FromC(const XLA_HloModuleConfig&)`
`xla::ShapeIndex`	（C index）	`ToC(const ShapeIndex&)`
`stream_executor::DeviceAddressBase`	`SE_DeviceAddressBase`	`ToC(const DeviceAddressBase&)`、`FromC(const SE_DeviceAddressBase&)`
`stream_executor::ScopedDeviceAddress<uint8>`	`SE_ScopedDeviceAddress`	`ToC(ScopedDeviceAddress<uchar>*)`
`stream_executor::DeviceAddressAllocator`	（allocate/free 回调对）	`ToC(DeviceAddressAllocator)` — `$_0::__invoke` 产生 `SE_ScopedDeviceAddress`，`$_1::__invoke` 释放 `SE_DeviceAddressBase*`

所有权规则

ToC 分配；Destroy 释放。Destroy 重载只存在于三个带堆拥有内部数据的类型：XLA_Shape（嵌套维度/元组子形状）、XLA_Literal（数据缓冲区）、XLA_ShapedBuffer（设备地址数组）。固定大小句柄（XLA_Layout、SE_DeviceAddressBase）没有 Destroy，因为除了调用方拥有的结构体本身外没有要释放的内容。重新实现者必须把每个变量大小类型的 ToC 与匹配的 Destroy 配对，并且绝不能对仅 POD 的句柄调用 Destroy。

特性 — 变量大小类型的 ToC 是输出参数形式 ToC(const xla::Shape&, XLA_Shape* out)，它填充调用方提供的结构体，而不是返回新结构体。随后基于指针的 Destroy(XLA_Shape*) 释放的是内部数据（维度数组、元组子形状），如果结构体本身位于调用方栈上，则不释放结构体本身。把“释放结构体”与“释放其内部数据”混淆会导致泄漏或双重释放。遵循重载签名：输出参数填充配对内部释放。

这个约定让名册函数能够接收和返回丰富数据，而无需共享 C++ ABI：TpuTransferManager_TransferLiteralToInfeed 接收宿主通过 ApiConverter::ToC 产生的 XLA_Literal*，插件在内部用它自己的静态链接 XLA 重建 xla::Literal；两个 xla::Literal 布局永远不需要匹配，因为只有扁平的 XLA_Literal 形式跨越了边界。

4. 名册地图

C-ABI 表面按其支撑的 SE 抽象分割为若干簇。计数是 IDA 在此构建中为每个前缀恢复出的 extern "C" 自由函数数量。每个簇都有一个专门的同级页面，清点其函数、签名和槽位偏移；本地图只给出一行角色说明和链接。

名册	C-ABI 数量	支撑	详情页
`TpuCompiler_*`	7	XLA→TPU 编译：`New`、`Compile`、`RunHloPasses`、`RunBackend`、`ShapeSize`、`DefaultDeviceShapeRepresentation`、`Free`	TpuCompiler 名册
`TpuExecutable_*`	9	已编译 executable 句柄：反序列化、指纹、查询 HLO modules / layouts / cost	TpuExecutable 名册
`TpuExecutor_*`	25	单设备运行时：`Allocate`、stream/event 创建、依赖连线、内存操作	TpuExecutor 名册
`TpuTransferManager_*`	19	宿主↔设备传输：literal 到/从设备、infeed/outfeed、缓冲区访问谓词	TpuTransferManager 名册
`TpuProgram_*`	18	序列化 program 对象：`New`、从 proto 反序列化、指纹、executable-info、释放	TpuProgram 名册
`TpuPlatform_*`	11	支撑 `TpuPlatform` 的 platform 生命周期 / 设备枚举	TpuPlatform 与 TpuNodeContext
`TpuTopology_*`	17	网格几何：chip 边界（X/Y/Z）、core 数量、host 数量、core 查找	TpuTopology 与 TpuCoreLocation
`TpuEmbeddingEngine_*`	15	SparseCore embedding：配置 host/memory、连接 hosts、dedup-data 计算	TpuEmbeddingEngine ABI
`TpuConfigurationApi_*`	8	运行时配置：编译缓存服务器地址、pod-state 查询、数组释放	TpuConfigurationApi

四个较小的簇与九个命名名册并列存在，并记录在上述页面内，而不是拥有独立页面：

名册	C-ABI 数量	支撑	记录于
`TpuCoreLocation_*`	4	单 core 坐标 / id（`TpuCoreLocationExternal`）	TpuTopology 与 TpuCoreLocation
`TpuNodeContext_*`	5	单节点驱动 context 启动	TpuPlatform 与 TpuNodeContext
`TpuMeshState_*`	3	分布式 mesh state 句柄	TpuPlatform 与 TpuNodeContext
`TpuCompile_*`	6	独立 compile 入口（不同于 `TpuCompiler_*` 类方法）	TpuCompiler 名册

名册调用如何流动

从开源 SE 调用到插件实现的路径在每个簇中都是一致的。以 TpuExecutor::Allocate 作为示例：

text

xla / PjRtClient
  └─ stream_executor::tpu::TpuExecutor::Allocate(...)        (host-side C++ shim, in this binary)
       └─ tbl = ExecutorApiFn()                               0x20819360  — get the singleton struct
       └─ tbl->slot[Allocate](se_handle, size, ...)           — call through the fn-ptr slot
              │  (the slot was populated at init to point at:)
              ▼
       TpuExecutor_Allocate                                   0xeab9120   — plugin-side C-ABI impl
              └─ (*se_handle->vtable[+136])(...)               — into libtpu's real TPU driver core
```text

shim 方法从不命名 `TpuExecutor_Allocate`；它索引一个槽位。该槽位在初始化时被填充为 `&TpuExecutor_Allocate`。C-ABI 函数接收不透明 `SE_StreamExecutor*`（此处在 `**a2` 读取）并分派进真实驱动 vtable。这个三跳形状，即*访问器 → 槽位 → C-ABI 实现 → 驱动*，是重新实现者必须复现的唯一模式；名册之间只是在槽位和句柄上有所不同。

> **注意 —** 宿主侧 `TpuExecutor::*` shim 和插件侧 `TpuExecutor_*` C-ABI 函数都被编译进了 `libtpu.so`，因为这里静态链接了 XLA。在干净拆分中（宿主进程通过 `dlopen` 加载插件），`TpuExecutor::*` shim 位于宿主，`TpuExecutor_*` 实现位于插件；只有 `*ApiFn` 结构体指针跨越边界。此二进制把两半折叠进同一镜像，因此同一地址空间同时持有两类名称。

---

## 5. 与 PJRT C-API 共存

此二进制中同时存在两个 C ABI。这里记录的 **TfTpu SE shim** 是*遗留*表面：它早于 PJRT，并以扁平 C 暴露完整的 StreamExecutor 对象模型（platform / executor / stream / transfer-manager / compiler）。**PJRT C-API**（通过版本化 vtable 到达的 `PJRT_*` 函数）是*现代*公共插件 ABI，也是外部框架实际加载 `libtpu.so` 的目标。它们不是位于同一层级并排的替代方案，PJRT 位于 StreamExecutor *之上*。TPU 的 `PJRT_Client` 由 `TpuClient` 实现，后者驱动 `TpuPlatform` 产生的 `TpuExecutor`，而 `TpuPlatform` 通过 `ExecutorApiFn()` 转发。因此单个 `PJRT_Client_Compile` 调用会向下经过 PJRT vtable，进入 SE 适配器，最终落到通过本页映射的这些 `*ApiFn` 表调用的 `TpuCompiler_*` / `TpuExecutor_*` C-ABI 调用。

所有权划分是：PJRT vtable、其 extension chain 和 `GetPjRtApi` 入口点归属于 [PJRT 概览](../pjrt/overview.md) 及其子页面。把 PJRT 向下桥接到此 shim 的 SE platform/executor 对象模型归属于 [StreamExecutor Platform 与 Executor 模型](../pjrt/stream-executor-host-interpreter.md)。本页只负责最底层，即 C-ABI 边界及其访问器/句柄约定。

> **特性 —** 面向*新框架集成*的重新实现者应针对 PJRT 编写，而绝不直接针对 TfTpu shim；shim 是 TPU PJRT 插件构建方式的内部细节。TfTpu shim 对重新实现*插件内部*或 TensorFlow 中的*遗留 stream-executor TPU 后端*的人重要，而不是对下游消费者重要。两个 ABI 共存，是因为即使最外层调用来自 PJRT，XLA 内部仍通过 StreamExecutor 驱动 TPU。

---

## 相关组件

| 名称 | 关系 |
|---|---|
| `stream_executor::tpu::ExecutorApiFn / OpsApiFn / ProfilerApiFn` | 位于 C-ABI 名册前端的三个单例访问器 |
| `TfTpu_ExecutorApiFn`（结构体类型） | 设备运行时名册通过其到达的函数指针表 |
| `ApiConverter::ToC / FromC / Destroy` | 丰富 C++ 类型与 `SE_*` / `XLA_*` 不透明句柄之间的编组 |
| `stream_executor::tpu::TpuExecutor` / `tensorflow::tpu::TpuPlatform` / `xla::TpuCompiler` | 通过 `*ApiFn` 槽位转发每个方法的宿主侧 SE shim |
| `TpuExecutor_* / TpuCompiler_* / TpuProgram_* / …` | 槽位指向的插件侧 C-ABI 自由函数 |
| PJRT C-API (`PJRT_*`) | 分层位于此遗留 SE shim 之上的现代公共插件 ABI |

## 交叉引用

- [TpuCompiler 名册](tpu-compiler-roster.md) — `TpuCompiler_*` / `TpuCompile_*` 编译 C 表面
- [TpuExecutable 名册](tpu-executable-roster.md) — `TpuExecutable_*` 已编译 executable 句柄操作
- [TpuExecutor 名册](tpu-executor-roster.md) — `TpuExecutor_*` 单设备运行时 + stream/event C 函数
- [TpuTransferManager 名册](tpu-transfer-manager.md) — `TpuTransferManager_*` 宿主↔设备传输 C ABI
- [TpuProgram 名册](tpu-program-roster.md) — `TpuProgram_*` program 对象 + 序列化 C ABI
- [TpuPlatform 与 TpuNodeContext](tpu-platform-and-topology.md) — `TpuPlatform_*` / `TpuNodeContext_*` / `TpuMeshState_*` platform + node-context C ABI
- [TpuTopology 与 TpuCoreLocation](tpu-topology.md) — `TpuTopology_*` / `TpuCoreLocation_*` 网格几何 C ABI
- [TpuEmbeddingEngine ABI](tpu-embedding-engine.md) — `TpuEmbeddingEngine_*` SparseCore embedding C 表面（通过 `OpsApiFn` 到达）
- [TpuConfigurationApi](tpu-configuration-api.md) — `TpuConfigurationApi_*` 运行时配置 C 入口点
- [TfTpu_Initialize 引导](../lifecycle/tftpu-initialize-bootstrap.md) — `*ApiFn` 函数指针结构体的一次性填充
- [StreamExecutor Platform 与 Executor 模型](../pjrt/stream-executor-host-interpreter.md) — 消费 `ExecutorApiFn` 的 SE 对象模型（实际运行中的 `*ApiFn` 槽位模式）
- [PJRT 概览](../pjrt/overview.md) — 与此遗留 SE shim 共存并位于其上的现代 PJRT C-API

TfTpu C-API Shim ​

摘要 ​

1. 为什么存在 C Shim ​

ABI 边界 ​

两半分别位于何处 ​

2. *ApiFn() 访问器模式 ​

单例函数指针结构体 ​

3. 不透明句柄约定 ​

SE_* / XLA_* / Tpu* 镜像 ​

所有权规则 ​

4. 名册地图 ​

名册调用如何流动 ​