TpuEmbeddingEngine ABI

本页所有地址均适用于 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）。IDA 恢复的 C 名称和反混淆后的 C++ 符号均逐字引用；.text VMA 等于文件偏移。其他版本会不同。

摘要

TpuEmbeddingEngine_* 是面向 SparseCore embedding-offload 子系统 的 C-ABI 函数组：它是主机侧入口，XLA/TensorFlow 通过它驱动设备端 embedding 流水线（大型查找表、稀疏 gather/scatter、分段 combiner 归约、去重），同时不在 .so 边界两侧共享任何单个 C++ 类型。其他 shim 名册通过 ExecutorApiFn 触达 stream_executor 设备运行时抽象，而本名册通过 OpsApiFn 访问器（0x10900e80）触达，这是专用于 embedding/SparseCore ops 的表面分区。本页记录的都是这条接缝：序列化的 TPUEmbeddingConfiguration proto 和扁平 *_Params 结构体如何进入插件，每个 C 调用分派到哪个 tensorflow::* / BarnaCoreManager / barna_core_util 核心函数，以及结果如何流回。那些函数最终驱动的设备端原语，即 minibatching 分解、分段扫描 combiner、去重 multiplicity，归属于 ../sparsecore/，此处只链接而不重新推导。

十五个 extern "C" 自由函数清晰地分成四个源文件，这个划分也对应子系统生命周期。配置（tpu_embedding_engine_configuration_ops_c_api.cc）负责拉起引擎：在核心间划分表、计算并汇集 HBM、配置并连接主机、完成初始化、探测初始化状态。加载/取回（tpu_embedding_engine_load_retrieve_ops_c_api.cc）通过 BarnaCore 驱动在设备和主机之间移动参数张量。XLA ops（tpu_embedding_engine_xla_ops_c_api.cc）完全不在设备上执行；它们各自构建一个 xla::XlaBuilder 计算，由主机拼接进自己的 HLO 图（接收 activations、发送 gradients、dedup-data 三元组）。入队（tpu_embedding_engine_enqueue_ops_c_api.cc）把一个稀疏输入 batch 送入引擎的每主机 batch 创建器。第十六个关注点是引擎的资源句柄，即小型的 TpuEmbeddingEngineState_* 三元组，它包装配置调用所操作的长生命周期 tensorflow::TpuEmbeddingEngineState C++ 对象。

有两种 C-ABI 形态反复出现，重实现者必须同时复现。每个函数都只接受一个参数：指向扁平 *_Params 结构体 的指针，主机填充该结构；字段按字节偏移读取（a1+16、a1+24，等等）作为输入，并通过间接 out 指针写回（**(a1+40) = operator new(...)）作为输出，末尾有一个被调用方通过 tsl::Set_TF_Status_from_Status 填充的 TF_Status* 槽位。并且每个返回状态的核心调用都使用 absl::Status 的“OK 是哨兵 (char*)&dword_0 + 1”惯用法：等于 OK 表示成功，否则把错误编组进 out-status 并 Unref 该 rep。这与 shim 概览 中建立的不透明句柄/*ApiFn 规则相同；本页只记录 embedding 函数组的槽位。

对重实现而言，契约是：

• *_Params 结构体约定：每个调用一个 in-pointer；输入按固定偏移读取；输出通过 operator new + memcpy 返回到主机拥有的 out 指针；末尾是 TF_Status* 槽位。
• proto/序列化配置摄入：大多数条目先从 {ptr,len} 槽位执行 stream_executor::tpu::DeserializeProto<TPUEmbeddingConfiguration, TpuSerializedProto>，再按解码后的 config 分派。
• 四族分派映射：每个 C 入口调用哪个 tensorflow::* / BarnaCoreManager::* / barna_core_util::* 核心函数，以及它是执行（configuration / load / enqueue）还是降低到 HLO（XLA ops）。
• 状态句柄：TpuEmbeddingEngineState_* 是配置族所修改的资源，区别于每次调用的 *_Params。

访问器	stream_executor::tpu::OpsApiFn() @ 0x10900e80（embedding/SparseCore 分区）
C-ABI 自由函数	15 个（TpuEmbeddingEngine_*）+ 3 个状态句柄函数（TpuEmbeddingEngineState_*）
地址跨度（engine）	0xf6a5b20 – 0xf769be0
调用约定	单个 *_Params* 参数；按偏移寻址的 in/out 字段；末尾 TF_Status*
配置摄入	DeserializeProto<tensorflow::tpu::TPUEmbeddingConfiguration, TpuSerializedProto>
状态惯用法	absl::Status OK 哨兵 (char*)&dword_0 + 1；错误 → tsl::Set_TF_Status_from_Status
核心驱动	platforms_deepsea::jellyfish::barna_core::BarnaCoreManager（load/retrieve、enqueue）
XLA 降低	通过 xla::XlaBuilder 调用 tensorflow::barna_core_util::*Computation
状态句柄类型	tensorflow::TpuEmbeddingEngineState（由 resource-mgr 支撑）
证据等级	可重实现级 / 经 IDA 反编译逐字节确认

范围 — 这个 C-API 最终驱动的设备端 embedding ops，即 minibatching 分解、segmented-add combiner、去重 multiplicity，归属于 ../sparsecore/（见 Embedding 小批处理、去重 Multiplicity）。BarnaCoreManager 所编程的 BarnaCore 标量引擎归属于 ../barnacore/。*ApiFn 访问器 / 不透明句柄模式归属于 shim 概览。本页只负责 TpuEmbeddingEngine_* 函数名册和 host→SparseCore offload 入口映射。

---

1. *_Params 调用约定

目的

每个 embedding C-ABI 函数都接受一个指向主机分配的参数结构体的指针，并返回 void（少数“返回值”的函数也只是把状态 rep 由函数体自身存入结构体）。调用点没有名称信息：主机填充结构体，通过某个 OpsApiFn 槽位调用，然后读回结构体的输出字段。结构体布局就是 ABI 契约：重实现者一旦移动字段偏移，就会让每个调用方静默损坏，因为没有其他机制可兜底。

算法

形态是统一的。TpuEmbeddingEngine_IsInitialized 是最小的完整例子：从两个输入槽读取序列化配置字符串，调用核心谓词，写回一个 bool，设置状态。

function TpuEmbeddingEngine_IsInitialized(Params* p):       // sub_f6a6ba0
    // optional VLOG(1) site, file tpu_embedding_engine_configuration_ops_c_api.cc:121
    len = p->config_len            // p+16
    buf = p->config_ptr            // p+24 — serialized TPUEmbeddingConfiguration bytes
    // small-string-optimization copy into a local std::string (len>22 => heap)
    status = tensorflow::IsTPUEmbeddingInitialized(buf, len, &out_bool)
    if status != OK_SENTINEL:                                // (char*)&dword_0 + 1
        // store into p->status (p+40), Unref old rep
        store_status(p->status_slot, status)
    *(*(p+32)) = out_bool          // **p->out_initialized = result
    if heap_allocated: free(local)
```text

> **NOTE —** OK 哨兵是字面指针值 `(char*)&dword_0 + 1`，也就是一个内联表示编码为“OK”的 `absl::Status`（低位已置位、无堆 rep）。本名册中的每个函数体都会把返回的 rep 与该常量比较；*相等*就是成功。若重实现把非空指针视为错误 rep，就会误读每一次成功调用。

输入侧反向使用 libc++ 的 small-string optimization：主机传入 `{len, ptr}`，被调用方在栈上重建 `std::string`，仅当 `len > 22` 时才堆分配。输出侧与之镜像：`std::string`/字节 blob 结果（`ExecutePartitioner`、`ConfigureMemory`、`CollateMemory`、`ConfigureHost`）通过把字节数写入一个 out 槽位，并把一份新 `operator new` 出来的拷贝写入另一个 out 槽位来返回：

```c
// output marshalling shared by ExecutePartitioner / ConfigureMemory / CollateMemory   (e.g. sub_f6a5b20)
**(p+32) = result.size                       // out: byte count
**(p+40) = operator new(result.size)         // out: host receives an owned copy
memcpy(**(p+40), result.data, result.size)   // caller frees later via TpuConfigurationApi free

GOTCHA — 输出缓冲区由插件通过 operator new 分配，并作为裸指针交给主机；主机拥有它，且必须通过匹配的数组释放入口（TpuConfigurationApi_* free 函数，见 TpuConfigurationApi）释放，而不能使用自己的分配器。跨接缝混用分配器会破坏堆。这也是 engine 名册没有自己的 _Free 的原因：释放集中处理。

注意事项

*_Params 字段偏移因函数而异（每个结构体都是定制的），但语法固定：开头是输入标量/{len,ptr} 对，然后是 out* 双指针，最后是 TF_Status*。末尾状态约定使这些函数可以返回 void：失败通过结构体内带返回，而不是通过返回值。

---

2. 配置族

目的

tpu_embedding_engine_configuration_ops_c_api.cc 中的七个入口组成引擎拉起序列。它们在语义上有状态：在 SparseCore 间划分 embedding 表、预留并布局 HBM、把 pod 中的主机连起来，并把引擎翻转到“initialized”。但 C 表面保持按调用无状态：每个入口摄入序列化的 TPUEmbeddingConfiguration（对 memory 调用还摄入 partitioner 的先前输出），并返回下一阶段的序列化结果。

入口点

host (XLA TPU embedding configuration ops)
  └─ OpsApiFn()[slot]  ── 0x10900e80
       └─ TpuEmbeddingEngine_ExecutePartitioner   0xf6a5b20  ── partition tables → partitions blob
       └─ TpuEmbeddingEngine_ConfigureMemory      0xf6a5d80  ── size HBM for one core
       └─ TpuEmbeddingEngine_CollateMemory        0xf6a5fa0  ── merge per-core HBM into pod layout
       └─ TpuEmbeddingEngine_ConfigureHost        0xf6a6340  ── per-host config blob
       └─ TpuEmbeddingEngine_ConnectHosts         0xf6a6660  ── wire pod hosts
       └─ TpuEmbeddingEngine_Finalize             0xf6a6960  ── commit; engine ready
       └─ TpuEmbeddingEngine_IsInitialized        0xf6a6ba0  ── bool probe
```text

### 算法

`ExecutePartitioner` 是规范的配置入口：反序列化 config proto，运行 partitioner，再把 partitions blob 编组返回：

```c
function TpuEmbeddingEngine_ExecutePartitioner(Params* p):   // sub_f6a5b20
    cfg = DeserializeProto<TPUEmbeddingConfiguration>(p+16)   // TpuSerializedProto {ptr,len}
    status = tensorflow::ExecuteTpuEmbeddingPartitioner(cfg, &out_blob)
    if status == OK_SENTINEL:
        **(p+32) = out_blob.size
        **(p+40) = operator new(out_blob.size); memcpy(...)   // partitions proto bytes
    else:
        Set_TF_Status_from_Status(p->status, status)          // p+48
    ~TPUEmbeddingConfiguration(cfg)

ConfigureMemory（0xf6a5d80）读取一个 core index（p+16）和一个序列化 partitions blob（p+24/+32），并调用 tensorflow::ConfigureTpuEmbeddingMemory(core, blob, len, &out)。CollateMemory（0xf6a5fa0）摄入一个逐核心 memory blob 的向量：它循环 p[+16] 次，把每个 {len,ptr} 以每元素 24 字节 stride 复制进 std::vector<std::string>，再调用 tensorflow::CollateTpuEmbeddingMemory(vec.data, vec.size, &out)，然后逐元素释放临时向量。ConfigureHost（0xf6a6340）以 host index 执行与 ConfigureMemory 相同的 string-in/blob-out 形态，并分派到 tensorflow::ConfigureTpuEmbeddingHost。ConnectHosts 分派到 tensorflow::ConnectTpuEmbeddingHosts，Finalize 分派到 tensorflow::FinalizeTpuEmbedding；二者都消费 config 且只报告状态（没有大型输出 blob）。

函数映射

函数	地址	大小	核心被调用方
TpuEmbeddingEngine_ExecutePartitioner	0xf6a5b20	363	tensorflow::ExecuteTpuEmbeddingPartitioner
TpuEmbeddingEngine_ConfigureMemory	0xf6a5d80	540	tensorflow::ConfigureTpuEmbeddingMemory
TpuEmbeddingEngine_CollateMemory	0xf6a5fa0	908	tensorflow::CollateTpuEmbeddingMemory
TpuEmbeddingEngine_ConfigureHost	0xf6a6340	774	tensorflow::ConfigureTpuEmbeddingHost
TpuEmbeddingEngine_ConnectHosts	0xf6a6660	738	tensorflow::ConnectTpuEmbeddingHosts
TpuEmbeddingEngine_Finalize	0xf6a6960	549	tensorflow::FinalizeTpuEmbedding
TpuEmbeddingEngine_IsInitialized	0xf6a6ba0	385	tensorflow::IsTPUEmbeddingInitialized

QUIRK — ConfigureMemory → CollateMemory 是 fan-in。ConfigureMemory 会对每个 SparseCore 调用一次，每次返回该核心的 HBM-layout blob；随后 CollateMemory 接受这些 blob 的整个数组，并把它们合并成一个 pod 范围布局。若重实现只用单个核心的输出调用 CollateMemory，就会生成忽略跨核心表分片的布局。CollateMemory 中每元素 24 字节的循环是在重建主机压平成 *_Params 结构体的 std::vector<std::string>。

注意事项

此构建中没有 TpuEmbeddingEngine_PartitionStateVariables 入口；partitioning 关注点是 ExecutePartitioner（表划分），它分派到 tensorflow::ExecuteTpuEmbeddingPartitioner。配置族不会通过这些 C 调用直接触碰 TpuEmbeddingEngineState 句柄；状态对象在主机侧通过 tensorflow::GetAndInitializeTpuEmbeddingEngineState(ResourceMgr*)（0xf78aec0）这一内部 C++ helper 单独解析，它不是扁平 C 名册的一部分。

---

3. 加载 / 取回族

目的

tpu_embedding_engine_load_retrieve_ops_c_api.cc 中的两个入口在主机和设备之间移动 embedding-table 参数张量。不同于配置族，它们调用实时的 BarnaCoreManager 驱动：它们要求引擎已初始化，并操作设备内存，而不是序列化 config blob。

入口点

TpuEmbeddingEngine_WriteParameters   0xf6a6d40
  └─ tensorflow::GetBarnaCoreManager(&mgr)                       ── resolve live driver
       └─ BarnaCoreManager::WriteParameters(mgr)                 ── host → device tables
TpuEmbeddingEngine_ReadParameters    0xf6a7160
  └─ tensorflow::GetBarnaCoreManager(&mgr)
       └─ BarnaCoreManager::ReadParameters(mgr)                  ── device → host tables
```text

### 算法

```c
function TpuEmbeddingEngine_WriteParameters(Params* p, TF_Status* status):   // sub_f6a6d40
    // build std::array<std::vector<absl::Span<float const>>, 8> from p
    //   outer dimension 8 == per-SparseCore fan-out; inner vector grows via
    //   2x reallocation (operator new(16*cap)) holding {ptr,len} float spans
    for core in 0..8:
        for slot in 0..p->count(p+64):
            span = { p[core][slot].data, p[core][slot].len }   // absl::Span<const float>
            push_back(spans[core], span)
    if GetBarnaCoreManager(&mgr) == OK and mgr != null:
        status = BarnaCoreManager::WriteParameters(mgr)         // pushes spans to device
    else if mgr == null:
        status = MakeErrorImpl<3>("TpuEmbeddingEngine not initialized.", file:43)
    Set_TF_Status_from_Status(status_out, status)
    ~array(spans)                                                // free inner vectors

ReadParameters（0xf6a7160，1000 字节）是对称的反向操作：解析 manager，调用 BarnaCoreManager::ReadParameters，把设备表编组回主机缓冲区。

函数映射

函数	地址	大小	核心被调用方
TpuEmbeddingEngine_WriteParameters	0xf6a6d40	750	BarnaCoreManager::WriteParameters
TpuEmbeddingEngine_ReadParameters	0xf6a7160	1000	BarnaCoreManager::ReadParameters

GOTCHA — 当 GetBarnaCoreManager 成功但返回空 manager 时，会触发显式的 "TpuEmbeddingEngine not initialized." 错误（字符串位于文件第 43 行）。若重实现只检查 GetBarnaCoreManager 的状态而不检查 manager 指针，就会解引用 null 并崩溃，而不是返回干净的 INVALID_ARGUMENT（MakeErrorImpl<3>）。

注意事项

std::array<..., 8> 的外层维度是逐 SparseCore fan-out（目标代际每个芯片八个核心）；每个核心的参数 span 被分别收集，因为设备端表按核心分片。内层 std::vector<absl::Span<const float>> 使用标准 libc++ 几何增长（反编译中的 2*cap / __throw_length_error 防护是 vector 重新分配，不是引擎逻辑）；重实现者可以忽略重分配算术，把它建模为 push_back。

---

4. XLA-Ops 族

目的

tpu_embedding_engine_xla_ops_c_api.cc 中的五个入口在类别上完全不同：它们发出 XLA HLO，并不在设备上运行。每个入口都构建一个 xla::XlaBuilder 计算，由主机的 XLA op-kernel 拼接进周围图中，使 embedding recv/send 和 dedup 操作成为 XLA 编译器可调度、可融合的一等 HLO。它们是 embedding 子系统与主 TPU 程序之间的桥梁。

入口点

TpuEmbeddingEngine_RecvActivationsComputation                   0xf767960
  └─ GetEmbeddingPartitionsProtoAndHbmBuffersConfig(...)        ── partitions + HBM layout
  └─ DeserializeProto<TPUEmbeddingConfiguration>(...)
  └─ xla::XlaBuilder::XlaBuilder(...)                           ── build HLO computation
TpuEmbeddingEngine_SendTPUEmbeddingGradientsComputation         0xf768d80
  └─ tensorflow::barna_core_util::LowerSendTPUEmbeddingGradientsComputation
TpuEmbeddingEngine_RecvTPUEmbeddingDeduplicationDataComputation 0xf7683e0
  └─ tensorflow::barna_core_util::LowerRecvTPUEmbeddingDeduplicationDataComputation
TpuEmbeddingEngine_DedupDataSizeComputation                     0xf7697e0
  └─ tensorflow::barna_core_util::DedupDataSizeComputation
TpuEmbeddingEngine_DedupDataTupleMaskComputation               0xf769be0
  └─ tensorflow::barna_core_util::DedupDataTupleMaskComputation
```text

### 算法

dedup-size 入口展示了共享前导：解析 embedding partitions proto 和 HBM-buffers config，解析拓扑，然后调用 `barna_core_util` lowering。

```c
function TpuEmbeddingEngine_DedupDataSizeComputation(Params* p):   // sub_f7697e0
    cfg = DeserializeProto<TPUEmbeddingConfiguration>(p+16)
    status = GetEmbeddingPartitionsProtoAndHbmBuffersConfig(
                 &partitions /*out, p+32*/, &hbm /*out, p+48*/)
    if status != OK: { Set_TF_Status(p+88, status); cleanup; return }
    if p->has_explicit_topology(p+72) == 0:
        topo = tfrt::GetGlobalTpuTopology()              // fall back to global
    else:
        topo = GetTpuTopology(p->serialized_topology)    // from p+72
    result = tensorflow::barna_core_util::DedupDataSizeComputation(
                 cfg, partitions, hbm, topo)             // file:324
    **(p+80) = result.size                               // out: dedup data size (int)
    // ~HbmBuffersConfig, ~EmbeddingPartitionsProto, ~TPUEmbeddingConfiguration

RecvActivationsComputation（0xf767960）执行相同前导，然后构造完整的 xla::XlaBuilder 并组装 recv-activations HLO（它会操作 xla::OpSharding 来放置结果），返回序列化计算。SendTPUEmbeddingGradientsComputation 和 RecvTPUEmbeddingDeduplicationDataComputation 分派到匹配的 barna_core_util::Lower*Computation builder；DedupDataTupleMaskComputation 分派到 barna_core_util::DedupDataTupleMaskComputation。

函数映射

函数	地址	大小	核心被调用方	发出
TpuEmbeddingEngine_RecvActivationsComputation	0xf767960	2463	xla::XlaBuilder recv-activations build	HLO computation
TpuEmbeddingEngine_RecvTPUEmbeddingDeduplicationDataComputation	0xf7683e0	2440	barna_core_util::LowerRecvTPUEmbeddingDeduplicationDataComputation	HLO computation
TpuEmbeddingEngine_SendTPUEmbeddingGradientsComputation	0xf768d80	2645	barna_core_util::LowerSendTPUEmbeddingGradientsComputation	HLO computation
TpuEmbeddingEngine_DedupDataSizeComputation	0xf7697e0	999	barna_core_util::DedupDataSizeComputation	size（int out）
TpuEmbeddingEngine_DedupDataTupleMaskComputation	0xf769be0	1325	barna_core_util::DedupDataTupleMaskComputation	tuple mask

QUIRK — *Computation 后缀是字面的：这些函数返回的是 HLO，不是设备结果。RecvActivationsComputation 实例化 xla::XlaBuilder 并发出 xla::OpSharding 就是明确信号：embedding lookup 被表达为主编译器调度的 XLA ops，因此 SparseCore recv/send 会融合进周围 step，而不是一次单独 dispatch。若重实现把它们建模为设备调用（像 load/retrieve 族那样），就会漏掉它们的输出是 XLA pipeline 仍需编译的子图。dedup 三元组（Size → TupleMask → recv DeduplicationData）是去重方案所需的主机侧元数据；设备端 multiplicity 数学在 去重 Multiplicity 中。

注意事项

拓扑解析是每个 XLA-op 入口共享的双分支：如果 *_Params 携带显式序列化拓扑（p+72 非零），则通过 GetTpuTopology 反序列化；否则引擎回退到进程全局 tfrt::GetGlobalTpuTopology()。全局路径会对 topology 取得共享引用（清理时 _InterlockedExchangeAdd64 递减引用计数），显式路径则堆包装一个新的 TpuTopology。二者都喂给同一个 barna_core_util lowering；重实现者必须保留两条路径，因为已配置 pod 运行和单主机运行会走不同路径。

---

5. EnqueueTensorBatch 与状态句柄

目的

最大的 engine 函数（EnqueueTensorBatch，4001 字节）和三函数状态句柄组不属于 configure/load/lower 族。EnqueueTensorBatch 把稀疏输入 minibatch 送入引擎的每主机 batch 创建器；TpuEmbeddingEngineState_* 三元组是配置族内部 C++ 侧操作的资源句柄。

算法

function TpuEmbeddingEngine_EnqueueTensorBatch(Params* p):       // sub_f6a9680
    // file tpu_embedding_engine_enqueue_ops_c_api.cc
    GetBarnaCoreManager(&mgr)
    creators = BarnaCoreManager::GetUserBatchCreators(mgr)       // per-host batch builders
    // pack the input tensor batch (indices/values/weights) into creators
    status = BarnaCoreManager::Run(mgr, ...)                     // drive minibatching
    Set_TF_Status_from_Status(p->status, status)
```text

`EnqueueTensorBatch` 是进入 **minibatching** 分解的主机入口：`BarnaCoreManager::GetUserBatchCreators` 返回逐主机对象，用于把 ragged input batch 转换为逐 SparseCore minibatch 格式；`BarnaCoreManager::Run` 驱动设备端 combine。分解本身见 [Embedding 小批处理](../sparsecore/embedding-minibatching.md)。

状态句柄有意保持简单：

```c
TpuEmbeddingEngineState_GetState(handle):   return *(void**)handle    // sub_f766fe0, 4 bytes
TpuEmbeddingEngineState_Create(...):        // sub_21389240 — allocate + init State
TpuEmbeddingEngineState_Free(...):          // sub_f766ea0 — destroy + free

函数映射

函数	地址	大小	角色
TpuEmbeddingEngine_EnqueueTensorBatch	0xf6a9680	4001	enqueue sparse batch → BarnaCoreManager::Run（minibatching）
TpuEmbeddingEngineState_Create	0x21389240	1295	分配并初始化 TpuEmbeddingEngineState
TpuEmbeddingEngineState_Free	0xf766ea0	292	析构并释放状态句柄
TpuEmbeddingEngineState_GetState	0xf766fe0	4	解引用句柄 → 内部状态指针

NOTE — TpuEmbeddingEngineState 是一个 tensorflow::TpuEmbeddingEngineState C++ 对象，在主机侧通过 TF ResourceMgr 触达（tensorflow::GetAndInitializeTpuEmbeddingEngineState，0xf78aec0；GetTpuEmbeddingEngineState，0xf78b1e0）。它的构造函数接受 tpu::TpuTopology（0xf961920）或 tpu::System（0xf9619e0），并在首次使用时运行 InitializeBarnaCoreManager（0xf961dc0，3204 字节），即从本引擎通向 BarnaCore 标量驱动的桥。这些 mangled tensorflow::TpuEmbeddingEngineState* 符号是内部 C++，不是扁平 C 名册的一部分；只有 Create / Free / GetState 跨越接缝。

注意事项

GetState 的 4 字节函数体（return *a1）确认该句柄是一个薄的 pointer-to-pointer 包装：主机持有的 C Tpu...State* 比真实 tensorflow::TpuEmbeddingEngineState 多一层间接。重实现应把该句柄复现为一个不透明的单指针结构，其第一个字就是状态指针；Create 分配两者，Free 拆毁两者，GetState 读取内部字。重活（BarnaCore manager 初始化）在 C++ 状态对象内部惰性发生，而不是在 Create 发生。

---

相关组件

名称	关系
stream_executor::tpu::OpsApiFn()	访问器（0x10900e80），其槽位持有这些 TpuEmbeddingEngine_* 指针
tensorflow::TpuEmbeddingEngineState	配置族所操作的、由 resource-mgr 支撑的 C++ 状态
platforms_deepsea::jellyfish::barna_core::BarnaCoreManager	load/retrieve 和 enqueue 背后的实时驱动
tensorflow::barna_core_util::*Computation	XLA-ops 族背后的 HLO lowering helper
stream_executor::tpu::DeserializeProto<TPUEmbeddingConfiguration, …>	大多数入口共享的序列化配置摄入
TpuConfigurationApi_*（array frees）	释放这些函数交回的 operator new 输出 blob

交叉引用

• TfTpu C-API Shim 概览 — *ApiFn 访问器 / 不透明句柄模式；本函数组通过 OpsApiFn 触达
• TpuExecutor 名册 — 通过 ExecutorApiFn 触达的设备运行时兄弟函数组
• TpuConfigurationApi — 释放 engine 输出 blob 的运行时配置 C 入口点
• SparseCore 概览 — 这些 C 调用所驱动的设备端 embedding engine
• Embedding 小批处理 — EnqueueTensorBatch 所喂入的 minibatch 分解
• 去重 Multiplicity — DedupData* XLA ops 背后的设备端数学
• BarnaCore 概览 — BarnaCoreManager 为 load/retrieve 和 enqueue 所编程的标量驱动