Skip to content

重构原始索引

此页面上的每个消息名称、字段计数和描述符 VA 均从 libtpu-0.0.40-cp314 轮中的 libtpu.solibtpu.so 描述符池中读取(BuildID md5 89edbbe81c5b328a958fe628a9f2207d,轮 0.0.40,未剥离)。字段编号是稳定的电线标识符;其他版本可能会添加字段,但不会重新编号。其他版本有所不同。

摘要

libtpu.so 将其整个 protobuf 类型 Universe 作为序列化 google.protobuf.FileDescriptorProto 记录嵌入到 protodesc_cold 部分(VA 0xbe8af30,大小 0x334180 = 3.20 MiB)。运行时在启动时重建这个描述符池,为每个生成的_pbz_…符号调用DescriptorPool::BuildFile();然后反射层按名称服务每条消息。完整的枚举 — 760 个文件记录、约 8 021 条消息、约 2 031 个枚举 — 均在同级 protodesc-cold-catalog 页面上编目。此页面是更窄、更高价值的索引:在这大约 760 个描述符中,哪些消息使该 wiki 逐个字段重建,以及每个重建在 哪个页面 上进行。

这里的“重构”意味着消息模式从描述符字节解码到重新实现级别——字段号→名称→标签→线路类型→嵌套类型结构→oneof分组——足以使读者可以重新生成.proto并往返线路格式。该索引与单纯的“编目”不同(同级页面按名称/大小/语法列出了所有 760 个索引,但不会对其中的大多数进行现场转储)。每行都带有描述符 VA,以便验证者可以使用 FileDescriptorProto.ParseFromString 重新解析源字节,并带有置信度标签,以便重新实现者知道要逐字信任哪些模式。

重建的集簇分为五个系列,每个系列由一个深层页面拥有:XLA HLO 序列化核心(前端 ↔ TPU 后端合约:HloModuleProto / HloInstructionProto / HloComputationProto 以及它们所依赖的 xla_data.proto 类型宇宙); 编译器配置原型(DebugOptions,290 个字段;TpuCompilationEnvironment,1121 个字段); TPU 可执行文件 + 拓扑原型(TpuCoreProgramProtoTpuChipPartsProto、共享内存定位器); 运行时/集体原型(Megascale 图、tpu_telemetryContinuationQueue 配置、环形路由缓存编解码器);和 探查器原型(XPlane 层次结构,xprof Task 记录)。下面分组的详细信息部分给出了每个系列的字段计数、描述符地址和所属页面。

对于使用该指数的读者,合约是:

  • 消息身份 — 完全限定的原始消息名称,因为它在池中显示为描述符字符串(可由 rg 通过二进制字符串进行验证)。
  • 字段计数此构建的描述符中声明的字段数量(不是您可能在其他地方找到的上游 .proto;计数是二进制事实,可能与公共 XLA 不一致)。
  • 所属页面 — 承载逐场重建的深层页面;该索引永远不会重复字段名册,它指向它。
  • 描述符 VA — 封闭 FileDescriptorProtoprotodesc_cold 地址,因此可以从源字节重新派生模式。
  • 置信度 - CERTAIN(描述符完全解码+针对所属页面和二进制文件进行交叉检查),HIGH(描述符解码,所属页面编写),MEDIUM(已解码但所属页面上的部分字段转储)。
描述符池protodesc_cold @ VA 0xbe8af30,大小 0x334180 (3.20 MiB)
池中的记录760 FileDescriptorProto; ~8 021 条消息,~2 031 个枚举
重建(此索引)跨越 5 个系列的约 30 条消息,约 14 个拥有页面
完整枚举protodesc-cold-catalog — 按名称/大小/语法列出的所有 760 个
标题计数DebugOptions 290 个字段 · TpuCompilationEnvironment 1121 个字段
重新解析路径FileDescriptorProto.ParseFromString(bytes_at_VA)

重建原始索引

此 wiki 已逐字段重建每条消息一行。 Fields 列是此构建的描述符计数; 所属页面是现场花名册所在的位置; 描述符 VAprotodesc_cold 中的封闭 FileDescriptorProto。 HLO 核心的描述符 VA(hlo.protoxla_data.protoxla.proto)来自文件记录映射;每条消息的 VA 是声明它的文件的 VA。

留言字段所属页面描述符 VA(文件)
xla.HloModuleProto25编译器/hlo 摄取0xc189a60 (hlo.proto)
xla.HloComputationProto8编译器/hlo 摄取0xc189a60 (hlo.proto)
xla.HloInstructionProto70编译器/hlo 摄取0xc189a60 (hlo.proto)
xla.HloScheduleProto1编译器/hlo 摄取0xc189a60 (hlo.proto)
xla.HloProto2编译器/hlo 摄取0xc189a60 (hlo.proto)
xla.ShapeProto5编译器/hlo 摄取0xc1b7e20(xla_data.proto)
xla.OpSharding14编译器/hlo 摄取0xc1b7e20(xla_data.proto)
xla.OpMetadata16编译器/hlo 摄取0xc1b7e20(xla_data.proto)
xla.HloModuleConfigProto42编译器/hlo 摄取0xc021470(xla.proto)
xla.DebugOptions290配置/调试选项-proto0xc021470(xla.proto)
xla::jellyfish.TpuCompilationEnvironment1121配置/tpu-编译-环境(tpu_compilation_environment.proto)
tpu.TpuCoreProgramProto11编译器/tpu程序序列化(tpu_core_program.proto)
tpu.TpuSequencerProgramProto5编译器/tpu程序序列化(tpu_sequencer_program.proto)
tpu.TpuChipPartsProto9目标/芯片零件-binarypb(tpu_chip_parts.proto)
tpu.TpuSharedMemoryOnChipProto2dma/连续队列(tpu_chip_parts_locators.proto)
TpuChipConfigProto.ContinuationQueue7+dma/连续队列(tpu_chip_config.proto)
…tpu_telemetry.CurrentCoreStateSummary7分析/tpu-遥测-原型(tpu_telemetry.proto)
…tpu_telemetry.AllCoreStateSummaries1分析/tpu-遥测-原型(tpu_telemetry.proto)
xla.megascale.runtime.MegaScaleInfoProto8超大规模/概述(megascale_info.proto)
xla.megascale.runtime.MegaScaleRuntimeError13超大规模/误差聚合器(megascale_status.proto)
xla.megascale.runtime.RapidEyeErrorDigestProto17超大规模/误差聚合器0xc169340 (rapideye_logging.proto)
CompressedToroidalRouteCache路由/环形路由缓存(route_cache.proto)
tensorflow.profiler.Task18分析/任务原型0xbe999a0(任务.proto)
tensorflow.profiler.XPlane6分析/xplane-xstat-traceme(xplane.proto)
tensorflow.profiler.XEvent4分析/跟踪 xevent 条目(xplane.proto)

注意 — 字段 列是*该消息上声明的顶级字段的计数,而不是文件的总数。 HloModuleProto直接声明25个字段; HloInstructionProto 是 70 个字段的广泛“所有操作属性的联合”;文件hlo.proto作为一个整体携带22条消息。文件总数位于 protodesc-cold-catalog 同级上。

明白了 — 这些计数是 此构建的二进制事实,其中一些与您在上游找到的公共 XLA .proto 不一致。 DebugOptions 这里有 290 个活动字段,最大字段数为 501(退役 GPU/CPU 标志有 211 个编号间隙)。驱动重新实现脱离描述符计数,而不是脱离上游 .proto


系列 1 — XLA HLO 串行化内核

JAX/TF/PyTorch 前端和 TPU JellyfishCompiler 之间的稳定合约。三个 FileDescriptorProto 记录承载整个模式,两个支持记录关闭类型图。依赖链是 xla_data.protohlo.protoxla.proto,因此这些记录中的 Universe 是独立的。

文件描述符 VA尺寸消息枚举角色
…/service/hlo.proto0xc189a6010 242225图脊柱:模块/计算/指令
…/xla_data.proto0xc1b7e2011 1435012类型 Universe:形状、布局、dtypes、分片
…/xla.proto0xc02147039 523130配置层:HloModuleConfigProtoDebugOptions
…/service/metrics.protoHloPassMetadata.kv_metrics 源码
…/hlo_profile_printer_data.proto每指令配置文件目录

图脊柱是 id-edge DAG,而不是树:HloModuleProto → 重复 HloComputationProto → 重复 HloInstructionProto,其中每个数据边都是同级指令列表中的 int64 operand_ids 引用,每个调用边都是模块计算中的 int64 called_computation_ids 引用列表。 root_id 每个计算都会命名输出。该索引表示在没有指针修复的情况下可以在序列化中幸存下来——这就是原型是一个干净的按索引的 DAG 的原因。

HloInstructionProto 是该系列中最广泛的消息:70 个字段,一个故意稀疏的“所有 op 属性的联合”,其中 string opcode = 2 选择哪个子集有意义(dot 读取 dot_dimension_numbers + precision_config;a custom-call 读取为 custom_call_target + backend_config + custom_call_api_version)。字段编号空间达到 99,其中退役属性有间隙。

QUIRK — HloOpcode 不是 原始枚举。操作码为string opcode = 2;对所有 760 个描述符的详尽扫描发现池中的任何位置都没有 xla.HloOpcode 枚举。操作码通过 C++ HloOpcodeString()StringToHloOpcode() 对序列化为小写文本助记符(“add”、“dot”、“卷积”、“fusion”、“all-reduce”)。这就是为什么原型在添加操作码的 XLA 版本之间向前/向后兼容的原因 — 新操作码只是一个新字符串,没有枚举号协调。从描述符池中构建数字操作码调度的重新实现者将找不到任何可调度的内容;助记词表是一个C++端表,而不是一个描述符。

注意 — HLO 核心重构位于 hlo 摄入,该页面拥有 StableHLO→HLO 转换和前端提交的 HLO 原型模式。id-graph 边缘模型、70 字段属性联合、分片方言(经典 OpSharding + Shardy) NamedShardingProto)和两层源来源(OpMetadata + StackFrameIndexProto)是从描述符池重建的。 xla_data.proto 类型 Universe(ShapeProtoOpShardingOpMetadata、维度编号消息)和 xla.proto 配置层 (HloModuleConfigProto) 使用相同的描述符记录,并与其一起进行描述。


系列 2 — 编译器配置原型

两个主标志消息控制整个编译管道,它们是索引中两个最高字段计数的重建。它们是兄弟姐妹:DebugOptions 是后端共享 (GPU/CPU/TPU) 标志集,与 HLO 模块一起跨越 PJRT 边界; TpuCompilationEnvironment 是 Jellyfish 后端实际使用的 TPU 特定主开关表。

留言字段最大字段#型号分解所属页面
xla.DebugOptions290501290 个实时数据、211 个编号间隙、17 个嵌套枚举、2 个地图条目消息配置/调试选项-proto
…TpuCompilationEnvironment11211218418 个布尔值、349 个消息、148 个 int64、74 个枚举、37 个字符串、34 个浮点、32 个 int32、14 个双精度、11 个 uint32、4 个 uint64配置/tpu-编译-环境

DebugOptions 住在 xla.proto (VA 0xc021470)。每个标量都是一个 proto3 optional,包裹在一个合成的单成员 oneof 中,给出了明确的 has-bit 存在 - 线上的 DebugOptions 准确记录了前端触摸的旋钮。 HLO-dump 和 HLO-pass-control 标志是管理该系列其余部分的序列化的部分:xla_dump_hlo_as_proto (113) → HloProtoxla_dump_hlo_snapshots (118) → HloSnapshotxla_dump_module_metadata (144) → HloModuleMetadataProtoxla_dump_latency_hiding_schedule (182) → ScheduleProtoxla_dump_full_hlo_config (381) → HloModuleConfigProtoxla_dump_hlo_unoptimized_snapshots (405) → HloUnoptimizedSnapshot

TpuCompilationEnvironment (TCE) 是目录中最大的单个原型(文件 137 692 字节)。它的 1121 个字段编号为 1…1218,其中有退休标志的间隙,并且至少 80% 包含在 AutoProto 中,因此每个字段都可以处于 AUTO / DISABLED / ENABLED 状态或带有键入的值。字段命名为子系统映射:xla_…xla_tpu_…xla_jf_…megascale_…xla_sc_…xla_tpu_sdc_checker_…xla_tpu_autofdo_…。这条消息是整个 TPU 编译器配置的表面区域;它的字段名册分为字段字典页面、偏移/默认页面和 AutoOr 解析器页面(所有页面都从所属页面交叉链接)。 DebugOptions 计数是权威的 290 个实时字段 — 请参阅 调试选项-原型 页面。

QUIRK — TCE 模式和标志表面是 相同​​的 1121 个名称,1:1:每个 TCE 字段也是注册的 absl::Flag。因此,重新实现者可以将原型字段名册和命令行/env 标志集视为一个命名空间,而不是两个。 TCE 是 editions 语法原型(proto2/proto3 的后继者); DebugOptions 是 proto3。两者均通过 CompilationEnvironmentsProto.environments 任意包装清单到达 RunHloPasses


系列 3 — TPU 可执行文件 + 拓扑原型

编译的“输出”端及其目标硬件描述。 TpuCoreProgramProto 是编译后的 TPU 可执行文件(JellyfishCompiler 的输出); TpuChipPartsProto 是每代硬件常数描述成本模型和 ISA 发射器读取。

留言字段角色所属页面
tpu.TpuCoreProgramProto11编译后的可执行文件; oneof core = TensorCore / BarnaCore / SparseCore编译器/tpu程序序列化
tpu.TpuSequencerProgramProto5每个定序器程序; oneof entry_point = 16 个每个芯片系列变体编译器/tpu程序序列化
tpu.TpuChipPartsProto9每个代号硬件常量(HBM/VMEM/SMEM、MXU 几何、时钟、DMA 颗粒)目标/芯片零件-binarypb
tpu.TpuSharedMemoryOnChipProto2共享内存定位器(类型+索引),由 ContinuationQueue 映射使用dma/连续队列

TpuCoreProgramProto.entry_point 确认支持的芯片系列完全是 {jellyfish、pufferfish (pxc)、viperfish (vxc)、ghostlite (gxc/glc)、6acc60406 (gxc/gfc)} 及其每核子程序(SCS = SparseCore Sequencer、TAC = Tile Access Core、TEC =平铺执行核心)。 tpu::l 序列化器将 oneof core 调度到判别式 *((_DWORD*)this + 22)(unsigned)(v20 - 5) <= 2 上,因此标签 5/6/7 是互斥的臂 — 请参阅所属页面的重建。

TpuChipPartsProto芯片零件-binarypb 上重建为嵌入式 <codename>_chip_parts.binarypb blob 的架构,运行时将其解析为 xla::jellyfish::Target 对象 - 一种数据驱动的 HAL,其中一个 C++ Target 类仅通过其加载的字节进行专门化。它的配套定位器文件 (tpu_chip_parts_locators.proto) 声明了 TpuCoreOnChipProtoTpuSharedMemoryOnChipProtoTpuSyncFlagRangeOnChipProtoTpuSegmentMemoryOnChipProto,从芯片部件架构和 ContinuationQueue 配置中引用。


系列 4 — 运行时/集体原型

驱动多芯片/多主机集合结构、设备状态快照、异步连续队列和环形路由缓存的线路格式。

留言字段角色所属页面
…MegaScaleInfoProto8每个集体操作元数据:TransferType、减少操作、不规则 + FP 压缩参数超大规模/概述
…MegaScaleRuntimeError13运行时错误记录:ErrorType、主机名、任务id、嵌入式运行时状态超大规模/误差聚合器
…RapidEyeErrorDigestProto17全船队根本原因摘要; 9 值 Cause 枚举,逐字嵌入第一个错误超大规模/误差聚合器
…tpu_telemetry.CurrentCoreStateSummary7每核实时状态:定序器信息、指纹、启动 ID、排队程序分析/tpu-遥测-原型
…tpu_telemetry.AllCoreStateSummaries1地图 int32 global_core_id → CurrentCoreStateSummary分析/tpu-遥测-原型
TpuChipConfigProto.ContinuationQueue7+异步连续队列配置:生产者标志、每核环窗口dma/连续队列
CompressedToroidalRouteCache烘焙路由缓存 blob,解压到每个链接的路由映射中路由/环形路由缓存

MegaScaleInfoProto 为每个跨切片集合键入元数据:TransferType ∈ {ALL_TO_ALL、ONE_TO_ONE、REDUCE_SCATTER、ALL_GATHER、ALL_REDUCE、BROADCAST、RAGGED_ALL_TO_ALL}、OperationType 缩减op、可选的参差不齐的参数和可选的 FP 压缩旋钮。它与 HloInstructionProto.channel_id + replica_group_list 处的 HLO 绑定,用于 TPU 后端在 DCN 上降低的集合。

RapidEyeErrorDigestProto 是最彻底重构的运行时原型:17 个实时字段、13 个嵌套消息类型、9 值 Cause 枚举,从 VA 0xc169340(4 460 字节)的 rapideye_logging.proto 记录中恢复。它逐字嵌入入站 MegaScaleRuntimeError(字段 11,first_recorded_error),因此摘要是自描述的。

ContinuationQueue 配置在 连续队列 上重建为静态记录(TpuChipConfigProto.ContinuationQueue,每个核心一个),它告诉主机运行时和设备生产者异步描述符环所在的位置 - 消息头上的生产者同步标志偏移、per_core 向量中的环区域和消费者标志。环形路由缓存 (CompressedToroidalRouteCache) 在 环形路由缓存 上的 MEDIUM 上重建:烘焙的 (src,dst) → route blob、其解压缩编解码器及其重复数据删除形式,由 GetDistanceFromCache 消耗,每对缓存未命中下降到实时环面指标。

明白了 — tpu_telemetry 是一个 时间点状态快照,在结构上与 XPlane 事件流相反 - 两者从不共享线路 blob。 AllCoreStateSummaries 是一瞬间捕获的每个核心一个 CurrentCoreStateSummary 的映射;探查器的 XSpace/XPlane/XLine/XEvent/XStat 是仅附加事件日志。重新实现者不得将它们统一起来;他们唯一的交汇点是都读取的 PC→HLO 元数据。


系列 5 — Profiler 原型

XPlane 跟踪层次结构和 xprof 每个工作线程元数据记录。这些是 XPlane 兼容的:带有内置元数据字典的 7 消息层次结构。

留言字段角色所属页面
tensorflow.profiler.Task18每个工作人员设备元数据:出处、主机 ID、时钟速率、配置文件窗口分析/任务原型
tensorflow.profiler.XPlane6跟踪平面:行 + 内部事件/统计元数据字典分析/xplane-xstat-traceme
tensorflow.profiler.XEvent4一个时间线事件:元数据 id、持续时间、统计信息、偏移量/发生次数之一分析/跟踪 xevent 条目

Task 是最精确锚定的探查器重建:18 个奇异标量字段的平面 proto3 消息,从 VA 0xbe999a0(路径 …/tensorboard_plugin_profile/protobuf/task.proto)处的描述符记录精确雕刻字节。在运行时,它是 Map<uint32 task_id, Task> 的“值”——每个工作序号一条记录——由 xprof::XlaJfProfileCheapOps(Task const&) @ 0xf2ca280 读取。 XPlane 层次结构(XSpaceXPlaneXLineXEventXStat,带有内部 XEventMetadata / XStatMetadata)在 xplane/traceme 和跟踪条目页面上重建;为其提供支持的每芯片 trace_entries.proto 系列在 跟踪点-主注册表 上单独编目。


验证来源

在该索引发布之前,根据二进制文件重新检查了每条消息名称和两个标题字段计数:

  • 消息名称存在 — 每个完全限定名称(HloModuleProtoHloInstructionProtoHloComputationProtoTpuCompilationEnvironmentDebugOptionsTpuChipPartsProtoTpuCoreProgramProtoTpuSequencerProgramProtoTpuSharedMemoryOnChipProtoMegaScaleInfoProtoMegaScaleRuntimeErrorContinuationQueueRouteCachetpu_telemetryCurrentCoreStateSummaryAllCoreStateSummaries)被确认作为描述符字符串出现在二进制文件提取的字符串。
  • 文件路径存在 - 确认存在源路径(service/hlo.protoxla_data.prototpu_chip_parts.prototpu_core_program.prototpu_sequencer_program.protomegascale_info.prototpu_telemetry.proto),锚定每个路径FileDescriptorProto 记录。
  • 标题计数DebugOptions = 290 和 TpuCompilationEnvironment = 1121 根据描述符解码和两个所属页面得到证实,它们独立地声明相同的数字。
  • 所属页面存在 — 索引表中链接的每个页面均已确认存在于 src/ 下。 HLO-core 系列行指向 hlo 摄入,该页面拥有前端 HLO 原型架构;配置行指向 config/ 字段字典页面;其余点位于现场转储每条消息的深层页面。

该索引做什么:它不会复制字段名册(那些位于所属页面上的字段名册),并且不会为未单独恢复其封闭文件 VA 的消息声明每消息描述符 VA — 这些行在 VA 列中携带文件名,而不是伪造的地址。空白的 VA 是一个诚实的间隙,而不是一个猜测的间隙。


交叉引用