Internal Wiki

主题

术语表 ​

本页中的所有名称和地址都适用于 libtpu-0.0.40-cp314 wheel 中的 libtpu.so（build-id 89edbbe81c5b328a958fe628a9f2207d；wheel/METADATA/__init__ 版本 0.0.40）。该二进制没有被 stripped：每个 C++ 符号都以 demangle 后的名称保留下来，因此这里说某个“术语有依据”，表示字面字符串或符号存在于二进制的名称/字符串表中。其他构建会不同。

摘要 ​

这是本 wiki 其余部分使用的每个缩写、代号和术语的规范定义来源。它存在的目的，是让其他页面无需顺带重新定义 MXU、SCS、Mosaic 或 gfc；它们会链接到这里，而这里为每个术语给出紧凑的一到三句定义，并指向拥有该术语的深度页面。读者如果在图中间落到一篇关于 SparseCore tile-execute sequencer 的页面，可以在几秒内在这里解析 TEC、Spmem 和 vex，然后跟随链接进入记录该机制的页面。

条目按领域分组：Silicon & Codenames、Compute Units、Memory Spaces、Compiler & IR、ABI & Runtime、Collectives & Network、Profiling 和 Reverse-Engineering Terms；每组内部按定义列表阅读。每个属于具体二进制事实的术语（codec namespace、sequencer ordinal、C-API struct）都会锚定到支撑它的符号、namespace 或 ELF section；属于外部事实的术语（营销名 “Trillium”、常规缩写 “PSUM”）或推断术语，会显式带有 (inferred) 或 (not in binary) 标签，而不是伪造锚点。

有两个结构性陷阱反复出现，值得在正文前先说明：单个 TPU generation 至少有七个名称，分散在三条数值不相交的整数轴上（TpuVersion、DeviceType、TpuVersionProto）；SparseCore sequencer enum 则存在两套相差一的编号（codec-template 与 proto/runtime）。二者都在下文定义，并在重新实现者最可能索引错误表的位置加以标记。

导航契约如下：

  • 每个有二进制依据的术语都会引用其依据：符号、namespace、ELF section 或拥有它的深度页面。
  • 不确定或外部术语会被标记为 (inferred)、(external) 或 (not in binary)；绝不会给它们伪造地址。
  • 两个反复出现的陷阱（三轴代号编号与两套编号的 sequencer enum）会在其所属位置给出提示。
Pluginlibtpu.so — Google TPU PJRT plugin，wheel 版本 0.0.40（build-id 89edbbe81c5b328a958fe628a9f2207d），约 745 MB
Companionsdk.so — proto/debug-event 分离体；见 Two-Binary Split
Codec 轴tpu::TpuVersion 0..5（dense）— 拥有 codename 绑定
Profiler 轴xprof::DeviceType 1..13（sparse）— 不同编号
ISA 层级LLO（Low-Level Optimizer）— 每个 bundle 的目标 IR
规范 codename 卡片Codename Cheat-Sheet

Silicon & Codenames ​

三套正交命名系统描述同一颗芯片：codec/ISA codename（嵌入 namespace 的两到三个字母标签）、fish codename（.rodata 中接近营销的内部名称）以及公开 Cloud 名称（v2…tpu7x）。codec codenames 在 asic_sw::driver::deepsea:: 下嵌套两层：先是家族标签（jxc/pxc/vxc/gxc），再是 fetch/load sub-core 标签。

NOTE — 本页中的 (N hits) 数字都是名称表 / sidecar 子串命中，也就是该 token 在符号名称索引中的每次出现；不是去重后的 nm 符号计数。它们是“二进制中有多少内容提到该术语”的数量级相关性信号，而不是不同函数或类的数量；单个类可通过其成员、vtable 和 type-info 字符串贡献数百次命中。若要查看去重后的总体数量（例如 RTTI _ZTI/_ZTV/_ZTS census），请使用 Forensics Overview，其中报告的是 nm 去重后的总数。不要把这些命中数相加，也不要把它们解读为基数。

术语定义
Jellyfish / jxc第一代 codec 家族；覆盖 Jellyfish（TpuVersion 0，“TPU v2”）和 Dragonfish（TpuVersion 1，“TPU v3”），二者共享一个没有 fetch/load 拆分的融合 codec。Namespace asic_sw::driver::deepsea::jxc::*（7,495 hits）。由 JXC Family 拥有。
Dragonfish“TPU v3” 的内部代号：TpuVersion 1，仍位于 jxc 家族内。字面量 dragonfish 位于 .rodata。
Pufferfish / pxc / pfc“TPU v4” 的 codec 家族；pxc::pfc 是 fetch sub-core（16,999 hits）。TpuVersion 2。由 PXC Family 拥有。
Puffylite / plcv4 级别的 lite 变体；load sub-core pxc::plc（2,174 hits）。它有自己的 profiler DeviceType（8），但没有自己的 TpuVersion：它折叠进 Pufferfish（TpuVersion 2）。
Viperfish / vxc / vfc“TPU v5” 的 codec 家族；vxc::vfc fetch sub-core（35,203 hits）。TpuVersion 3。由 VXC Family 拥有。
Viperlite / vlcv5 级别的 lite 变体（Viperfish-lite）；load sub-core vxc::vlc（4,223 hits）。DeviceType 11，折叠进 Viperfish（TpuVersion 3）。
Ghostlite / glcgxc 家族，load sub-core gxc::glc（126,927 hits）。TpuVersion 4，DeviceType 13，Cloud v6e。提供一个具名 codec 类 tpu::TpuCodecGhostlite。
gfc / 6acc60406 / Ironwoodgxc 家族，fetch sub-core gxc::gfc（126,696 hits）。TpuVersion 5，DeviceType 12，显示字符串 “TPU7x”。二进制中该芯片唯一的内部名称是混淆标签 6acc60406；“Ironwood”（v7x 营销名）和 “Ghostfish” 不在二进制中（外部名称）。不要与 “Trillium” 混淆，后者是前一代 v6e/Ghostlite。
TpuVersion内部 dense 0..5 codec/HAL 轴；TpuVersionToString 索引位于 off_22011BF0 的 6 指针表。这是 codec、HAL 和 ISA emitter switch 所用的轴。
DeviceTypeProfiler 的 sparse 1..13 轴（xprof::DeviceType）；由 DeviceTypeFromDeviceIdentifiers 按 PCI tuple 赋值。八个 TPU silicon 家族落在 {3,5,7,8,10,11,12,13} 上。
TpuVersionProtoprotobuf wire enum 1..6；internal = proto − 1。这就是为什么嵌入的 6acc60406_chip_parts.binarypb 对内部 TpuVersion 5 携带 version = 6。
HALHardware Abstraction Layer：TpuHal{Jxc,Pxc,Vxc}HardwareFactory 类。注意 gxc（Ghostlite + gfc）注册进共享的 Vxc factory；不存在 TpuHalGxc。
BarnaCore / bcs / brn早期 jxc 家族中存在的独立 scalar sub-core（jxc::bcs、jxc::brn）；一种不同于 TensorCore 的 32 字节 bundle scalar ISA。由 BarnaCore 拥有。

GOTCHA — TpuVersion ≠ DeviceType。 这是两个由两个子系统维护的 enum，从未意图对齐。TpuVersion 4 是 Ghostlite；DeviceType 4 不是任何 TPU generation。DeviceType 在顶端的编号也按时间倒置：Ghostlite（更旧，v6e）= 13，6acc60406（更新，v7x）= 12。若重新实现者用 compiler TpuVersion ordinal 去索引 profiler 的 per-device 表，会读到错误行。请通过 Codename Cheat-Sheet 解析。

NOTE — “Trillium”、“Ironwood”、“Ghostfish” 不在二进制中。 这些字面字符串出现次数为零。gxc::gfc 缩写可能代表类似 “Ghostfish” 的鱼类名称，但该名称不存在；应改为引用 6acc60406（内部）或 TPU7x（显示）。Trillium = v6e = Ghostlite 这个身份关系是正确的，但它是外部（Cloud 文档）事实。

Compute Units ​

TensorCore 是 dense-math 引擎；SparseCore 是面向 embedding/gather workload 的独立协处理器，各自拥有 sequencer-typed ISA。这里的缩写命名 bundle slot 所面向的每 slot 功能单元。

术语定义
TensorCore / TCdense matrix/vector core：驱动 MXU、VPU 和 XLU 的主要计算引擎。TensorCore 在名称表中出现 447,417 次；sequencer-type ordinal 0（codec-template）。横跨 ISA Overview。
MXUMatrix-multiply Unit：执行 dense dot/conv contraction 的 systolic array。符号家族 …::Mxu（476 hits）。由 MXU Slot 拥有；分配见 MXU Assignment Bin-Packer。
VPUVector Processing Unit：用于 elementwise vector ops 的 SIMD lane 引擎。…::Vpu（522 hits）。由 VPU Slot 拥有。
XLUCross-Lane Unit：lane-shuffle/reduction 单元（CrossLane*、…::Xlu，442 hits），用于在 VPU lane 之间移动数据（transpose、reduce、broadcast）。由 XLU Op Roster 拥有。
EUPExtended Unary Processor / transcendental unit：计算 exp、log、rsqrt 等。符号家族 …::Eup（1,025 hits）。由 EUP Transcendental Slot 拥有。
SPUScalar Processing Unit：每个 bundle 的 scalar ALU/address 引擎。SPU 出现（20 name hits）；即 bundle 的 scalar slot。由 SPU Scalar Slot 拥有。
SparseCore / SCembedding/sparse 协处理器；namespace SparseCore*（127,401 hits）。一套独立 ISA，带三种 sequencer type（SCS/TAC/TEC）。由 SparseCore Architecture 拥有。
SCSSparseCore Scalar sequencer：control/scalar 引擎。Codec-template sequencer ordinal 3（proto/runtime 4）。符号家族 SparseCoreScs*（6,425 hits）。由 SCS Engine 拥有。
TACSparseCore Tile-Access sequencer：驱动 gather/scatter tile addressing。Ordinal 4（proto 5）。SparseCoreTac*（5,966 hits）。由 TAC Engine 拥有。在 gfc/v7x 上被移除。
TECSparseCore Tile-Execute sequencer：运行每 tile compute。Ordinal 5（proto 6）。SparseCoreTec*（72,530 hits）。由 TEC Engine 拥有。
VEXSparseCore 内部的 Vector-Extended 单元：SparseCore vector datapath（mask/dest-port/sub-opcode 编码）。由 VectorExtended (vex) 拥有。

GOTCHA — 6acc60406（gfc/v7x）只提供 SCS + TEC，没有 TAC。 二进制中已确认：gfc::isa::SparseCoreScs*（1,139 hits）和 gfc::isa::SparseCoreTec*（17,193 hits）存在，但 gfc::isa::SparseCoreTac* 零命中。Viperfish（vfc）和 Ghostlite（glc）带有全部三个 sequencer；v7x 移除了 tile-access 引擎。若重新实现假设所有带 SparseCore 的 generation 都有统一的 SparseCore 三件套，就会发出 v7x 硬件没有 sequencer 可执行的 TAC codec。见 GetSequencerType。

GOTCHA — SparseCore sequencer enum 有两套相差一的编号。 codec template 实例化 {SCS=3, TAC=4, TEC=5}；proto/runtime 形式则高一位，为 {SCS=4, TAC=5, TEC=6}。TpuSequencerTypeFromProto 是字面上的 internal = proto − 1 switch。混用二者会静默编码到错误引擎（或编码到保留 slot 2/3）。见 Sequencer Ops Per Gen。

Memory Spaces ​

TPU 程序会寻址几个不同且不一致的 memory space，每个都有自己的 allocator 和 addressing 规则。space 是 LLO memory-space enum 的一部分（见 Memory-Space Enum）；下列名称来自二进制自身。

术语定义
HBMHigh-Bandwidth Memory：片外 DRAM，最大层级；保存程序输入/输出和 spilled state。符号家族 …Hbm…（27,346 hits）。由 HBM Allocator 拥有。
VMEMVector Memory：片上 scratchpad，用于供给 VPU/MXU 的 vector tile。…Vmem…（8,284 hits）。由 VMEM Allocator 拥有。
SMEMScalar Memory：片上 scalar scratchpad / SPU 的 register window。…Smem…（7,319 hits）。由 SMEM Scalar Memory 拥有。
CMEMCommon Memory：共享片上池（cross-core / staging）。…Cmem…（1,344 hits）。由 CMEM Pool 拥有。
SpmemSparseCore 私有 scratch memory（Spmem，13,122 hits）：AllocateDefaultSpmemWindow 为 SparseCore lowering 切出的缓冲空间。不同于 VMEM/SMEM。记录在 SparseCore 下。
PSUMPartial-Sum accumulator memory：MXU 的结果累加 bank。缩写 PSUM 不在二进制中；有依据的名称是 PartialSum*（979 name hits）。(external acronym)：grep 时使用 PartialSum。
SFLAGSync-Flag memory：用于 cross-engine synchronization 的小型 flag/semaphore 空间。有依据的形式是 SyncFlag*（9,882 hits）/ Sflag（1,158 hits）。由 SFLAG Protocol 拥有。
MRBMatrix-Result Buffer：write-back 前保存 MXU 输出的 FIFO/buffer；符号家族 …::Mrb…（82 hits）。由 MRB Chain Allocator 和 MRB FIFO/MSR Placement 拥有。

NOTE — PSUM 是约定，PartialSum 才是二进制。 wiki 使用 “PSUM”，因为它是 MXU accumulator bank 的标准 TPU 术语，但字面字符串 PSUM 不出现；每个二进制锚点都指向 PartialSum。不要搜索 PSUM 后就断定该功能不存在。

Compiler & IR ​

compiler 摄入高层 graph（HLO / StableHLO），经由 MLIR dialect（MHLO、TPU dialect、Mosaic）lower 到每个 bundle 的 LLO 目标 IR，随后调度 bundle 并分配 memory。下列名称是 compiler 页面读者会遇到的 dialect/pass 术语。

术语定义
HLOHigh-Level Optimizer IR：XLA 的经典 op-graph IR；compiler 的摄入形式。由 HLO Ingestion 拥有。
MHLOHLO 的 MLIR dialect 镜像：TPU lowering pipeline 的入口 MLIR dialect。见 MHLO→XTile/TPU Lowering。
StableHLO用于 JAX native serialization 的版本稳定 MLIR HLO dialect；二进制带有 StableHLO module op-doc string（71 string hits）。它是 MHLO 之前的可移植 interchange form。
MosaicTPU tiling/vector-layout MLIR 框架，将 tensor ops 映射到 VMEM tile 和 lane layout（Mosaic*，514 hits）。由 Mosaic Overview 拥有；layout inference 见 Mosaic Layout Inference。
TPU dialect位于 Mosaic 之下的 TPU-native ops 的 MLIR dialect，通过 ODS-generated pattern lowered 到 LLO。见 TPU Dialect and Ops 和 TPU→LLO ODS。
LLOLow-Level Optimizer IR：ISA emitter 消费的每 bundle 目标 IR；opcode enum 和 proto 形式存在于 Llo*（25,333 hits）。由 LLO Opcode Enum 拥有。
bundle一个 VLIW instruction word：由 per-unit slot（MXU/VPU/XLU/EUP/SPU/sequencer/memory）组成的固定宽度 pack。宽度按 generation 而定（41-bit JF、51-bit PF、64-bit VF 等）。由 Bundle Model Overview 拥有。
MSAMemory-Space Assignment：在 HBM/VMEM 等之间放置 buffer 并预留 segment 的 pass。有依据的形式是 MemorySpaceAssignment*（544 hits）。由 MSA Overview 拥有。
LatencyHidingScheduler / LHS将 compute 与 DMA/memory latency 重叠的 bundle scheduler（LatencyHiding*，363 hits）。由 Latency-Hiding Scheduler Core 拥有。
TpuCodec每个 TpuVersion 的 codec 对象，拥有 ISA encode/decode；由 6-case switch tpu::TpuCodec::Create(TpuVersion) 构建。符号家族 TpuCodec*（128 hits）。实例化 per-sequencer codec（以及 SCS/TAC/TEC ordinal）的 codec-template 挂在这里。
SPMD / shardingSingle-Program-Multiple-Data partitioning：把程序拆分到多个 core 的 auto-sharding pass；SpmdPartitioner*（596 hits）、ShardingPropagation*（143 hits）。由 Auto-Sharding SPMD 和 Sharding Propagation 拥有。
autotune记忆化最佳已发现 kernel config 的 cache（Autotune*，610 hits；autotune_results，8 hits）。Cache mode 为 READ/UPDATE/UNSPECIFIED；失效处理由用户负责。见 Fusion Cost Model。
SymbolicTile / IndexingMaptiling cost-model primitive：SymbolicTile 是一个 IndexingMap 三元组（offset-map、size-map、stride-map），描述 tile 如何映射到 tensor（SymbolicTile* 272，IndexingMap* 8,014 hits）。供 loop-tiling 和 fusion pass 使用。
addrspacecastMLIR/LLVM address-space cast op（addrspacecast*，1,578 hits）：在 memory space 之间重新给 pointer 定型的 lowering（尤其是 SparseCore fat-pointer space）。见 AddrSpaceCast ISel 和 Fat Pointers AS7/8/9。
walrus有时附加到 pass-pipeline driver 的术语，但在此二进制中（名称表和字符串表，不区分大小写）零次出现。(not in binary)：不要把任何主张锚定到它；pass-pipeline driver 的依据是普通的 xla::HloPassPipeline。如果某项功能曾被归因于 “walrus”，引用前应从有依据的符号重新推导。

ABI & Runtime ​

plugin 的外部表面是 PJRT C-API，分层构建在较旧的 TfTpu C-API 之上。Runtime 通过 infeed/outfeed 在设备内外移动数据，并在 stream 上异步执行。

术语定义
PJRTPortable JAX/XLA Runtime：framework 调用的 C-ABI plugin contract。plugin 的 struct family PJRT_Client / PJRT_Buffer / PJRT_Executable / PJRT_Device / PJRT_Event 全部存在。由 PJRT Overview 拥有。
PJRT_Api / GetPjrtApi顶层 dispatch vtable struct 及其 entry-point getter（GetPjrtApi，1 hit），framework 加载它以到达每个 PJRT 函数。在 API Vtable Reconstruction 中重建。
TfTpu C-API较旧的 TensorFlow-TPU C-API（TfTpu_Initialize 等，117 个 TfTpu hits），PJRT 分层构建其上；它仍是 executable execution 的底层基底。见 Stream-Executor / PJRT Adapter。
StreamExecutorXLA device-abstraction layer（StreamExecutor*，787 hits），PJRT 对其进行包装；管理 TPU 上的 stream、allocation 和 async event。见 Stream-Executor / Host Interpreter。
buffer lifecyclePJRT device-buffer 路径：BufferFromHostBuffer（upload）→ DonateWithControlDependency（alias-into-execute）→ Execute → Delete。外部 lifetime 通过 Increase/DecreaseExternalReferenceCount 引用计数。由 Buffer and Memory 拥有。
AsyncValue / ReadyEventasync-completion primitive：PJRT_Event（Await/OnReady/IsReady）包装一个 AsyncValue（5,088 hits）；buffer 的 ReadyEvent 表示其内容有效。由 Events and Async 拥有。
custom-call手写 kernel 的 tpu_custom_call HLO escape hatch（CustomCall*，4,319 hits）。plugin 没有 user FFI extension；custom op 只通过 tpu_custom_call registration 路由。见 Custom-Call Lowering。
DmaMap / cross-host receivePJRT Client-extension 表面（DmaMap/DmaUnmap，584 hits），用于 pin host memory 以进行直接 cross-host buffer transfer。由 DMA and Cross-Host Recv 拥有。
infeed / outfeedhost↔device streaming channel：infeed 将输入推送到 device queue，outfeed 排出输出（Infeed* 4,025 / Outfeed* 3,323 name hits）。由 Infeed/Outfeed 拥有。
ICIInter-Chip Interconnect：pod 内 chip-to-chip fabric（ICI，1,527 string hits）。在 slice/pod 内承载 collective traffic。由 ICI Overview 拥有。
MegascaleICI 之上的 cross-host control plane（xla::megascale::runtime::*）；通过 DCN 协调 multi-slice/multi-pod execution。由 Megascale Overview 拥有。
trace codecprofiler 组件，用于序列化 per-core trace entry；codec alternative 由原始 PCI identity 选择（不是由 TpuVersion/DeviceType 选择）。见 Riegeli Trace Container。

Collectives & Network ​

Collective communication（all-reduce、all-gather、reduce-scatter）在 pod 内通过 ICI 运行，在 host 间通过 DCN 运行。下列 topology 术语描述 collective 被调度到的 fabric。

术语定义
DCNData-Center Network：ICI 之上的 inter-host network（DCN，472 string hits）；DCNTopology 是其 topology proto。Megascale 通过它协调 transport。
superpod / podtopology 层级：chip 组成 pod（toroidal ICI mesh），pod 组成 superpod（HierarchicalTopology）。3D-Superpod resiliency mode 可容忍 sub-cube ICI link fault。见 Routing。
Megacore fusioncollective-lowering optimization，将 collective 与一颗芯片两个 core 上相邻 compute 融合。由 Megacore Fusion 拥有。
polarityICI link-assignment 属性（polarity state machine 在 topology bring-up 期间分配 link direction）。见 Link Bring-up。

Profiling ​

profiler 发出 Google XProf schema：一个 proto-defined event tree（XSpace → XPlane → XEvent + XStat），带有 per-core trace dispatcher，每种 sub-core type 一个。下列术语命名该 schema。

术语定义
XPlaneXProf schema 中的 timeline plane：每个 device 或 host 得到一个 plane（XPlane*，1,679 name hits）。由 XPlane / XStat / TraceMe 拥有。
XEventXPlane 上的单个 timed event（XEvent*，377 hits）：一个带 start/duration 和 metadata id 的 span。
XStat附着到 XEvent 或 XPlane 上的 typed key/value statistic（XStat*，244 hits）。
XSpace保存一次 capture 中所有 XPlane 的顶层容器（XSpace*，332 hits）。
TraceMe发出 scoped region 的 XEvent 的 instrumentation primitive（TraceMe*，1,121 hits）。
CoreDispatcherper-sub-core trace-entry decoder；二进制为每个 codec family（gfc、glc、jxc、pxc、vfc、vlc）实例化一个。由 Trace Entries Coder 拥有。
TpuProfiler ABIC-ABI profiler 表面（TpuProfiler_Create/Start/Stop/CollectData）。由 TPU Profiler ABI 拥有。

Reverse-Engineering Terms ​

这些是 wiki 用来描述事实如何从二进制恢复出来的 static-analysis 术语，而不是 TPU 本身的事实。

术语定义
vtableC++ virtual-method dispatch table；以 _ZTV… 符号发出（39,244 hits）。因为二进制未 stripped，vtable 地址可以解析到其所属 class。见 RTTI/Vtable Census。
RTTI / typeinfoRun-Time Type Information：编译器为每个 polymorphic class 发出的 _ZTI… type-info record（60,457 hits）和 typeinfo name string（141,445 hits）；这是为原本匿名的对象命名的主要杠杆。见 RTTI/Vtable Census。
.lrodata / .lbsslarge-code-model ELF section（二者都存在于此二进制中）。.lrodata 保存大型只读数据（codename table、proto descriptor），.lbss 保存大型 zero-init data；二者使用 64-bit relocation 寻址，而不是 ±2 GB small model。见 ELF Anatomy。
FLIRTFast Library Identification and Recognition Technology：IDA 的 library-function fingerprinting，用于标记静态链接的 Abseil/protobuf/tcmalloc/dnnl 代码，使其能从 TPU-specific analysis 中排除。嵌入库编目见 Embedded-Library Atlas。
sidecardisassembly 的 machine-readable export（_names.json、_strings.json、_segments.json、callgraph 和 ctree 文件），wiki grep 它们而不是重新运行 decompiler。本页每个计数都是 sidecar hit。
dispatch table编译出的 jump/pointer table（switch 或 per-TpuVersion function-pointer array），二进制在运行时索引它；per-generation 表是 codename derivation 的主干。见 Dispatch-Table Taxonomy 和 Per-Gen Function Dispatcher。
two-binary splitwheel 提供 libtpu.so（plugin）和 companion sdk.so（proto/debug-event definition）；许多 libtpu.sdk.* 符号位于后者。见 Two-Binary Split。

交叉引用 ​