TpuCompilationEnvironment
本页所有符号、地址和 struct 偏移均适用于
libtpu-0.0.40-cp314wheel 中的libtpu.so(构建libtpu_lts_20260413_b_RC00,build-id md589edbbe81c5b328a958fe628a9f2207d,781,691,048 字节,ELF x86-64 DYN,未 stripped)。字段名、proto 类型名和 mangled C++ 符号均逐字引用自二进制。其他版本会不同。
摘要
xla::jellyfish::TpuCompilationEnvironment(TCE)是 TPU 编译器的主配置 proto,是一个包含 1121 个字段的单一 message,保存 Jellyfish 后端读取的每个逐编译旋钮。它是后端共享 xla::DebugOptions 在 TPU 侧的伴生物:DebugOptions 携带约 290 个通用/dump/GPU/CPU 字段,这些字段随 HLO module 穿过 PJRT 边界,并且在 TPU 上大多惰性;而 TCE 携带 TPU 编译器实际消费的 1121 个字段,包括 scheduler tunables、fusion gates、MSA memory-space ratios、SparseCore offloading switches、collective combiner thresholds、BarnaCore embedding controls 以及 debug/trace toggles。它的定义性属性是:每个 TCE 字段也都是一个已注册的 absl::Flag,因此 proto schema 和 flag surface 是相同的 1121 个名称,1:1。
重新实现者的参照框架是 XLA 自身的 CompilationEnvironment 机制。TCE 是一个具体的 proto2::Message 子类,注册进逐 module 的 xla::CompilationEnvironments bag,并且按 C++ 类型取回,即 CompilationEnvironments::GetMutableEnv<TpuCompilationEnvironment>(@ 0xe6de1e0),而不是通过全局 singleton。默认实例通过遍历 proto descriptor 并复制每个字段注册的 absl::Flag 默认值来构建;这就是 CreateDefaultTpuCompEnv @ 0x1d73dfa0(逐字段构造器)和 GetTpuCompEnvWithDefaultValues @ 0x1d73f100(once-guarded 缓存默认值)。在这些 flag 默认值之上,一个按 TpuVersion 的 MSA overlay 会重写一组字段,因此 JAX 用户看到的值是 flag-default ⊕ AUTO-resolution-polarity ⊕ per-version-overlay。
本页是 TCE 的结构/地图:proto 是什么、如何构造、1121 个字段聚集到哪些子系统区域、关键嵌套 sub-message,以及 AutoOr<T> AUTO-vs-explicit wrapper 如何应用。它刻意不复现字段列表;完整 field#→name 字典拆分在 tce-field-dictionary-a.md/tce-field-dictionary-b.md,field→offset→default 参考在 tce-field-offsets-defaults.md,AUTO resolver 内部机制在 autoproto-autoor-resolution.md。本页负责分类法、sub-message 地图和构造路径。
对于重新实现,契约是:
- proto 身份和注册 — TCE 是一个
proto2::Message子类,注册为xla::CompilationEnvironment,从 HloModule 按 C++ 类型获取,而不是一个全局 config struct。 - 构造路径 —
CreateDefaultTpuCompEnv如何遍历 descriptor 并复制每个字段的absl::Flagdefault union 来构建默认值,以及缓存默认值如何只实体化一次。 - 字段组分类法 — 1121 个字段聚集到哪些子系统区域(让重新实现者知道某个旋钮位于哪里),以及界定 schema 的 proto 类型统计。
- sub-message + AutoProto 地图 — 类型化 sub-message(
RangeSpecProto、12 个 AutoProto message arm、嵌套TpuCustomCallMemorySpaceSpecpolicy)以及AutoOr<T>三态适用的位置。
| Message | xla::jellyfish::TpuCompilationEnvironment(package xla.jellyfish,proto2) |
| 字段数量 | 1121 个存活字段(最大 field# 1218,gap 已 tombstoned) |
sizeof | 0x15E8 (5608 B);字段数据区域 +0xA8 .. +0x15E0 |
| Parse table | _table_ @ 0x21cfa9e0(TcParseTableBase + 1121-entry FieldEntry 数组 @ +0x370) |
| 默认实例 | TpuCompilationEnvironment_globals_ @ 0x227b87e0 |
| Class data | TpuCompilationEnvironment_class_data_ @ 0x223c96a0(sizeof 0x15e8 at +0x20) |
| 构造默认值 | CreateDefaultTpuCompEnv @ 0x1d73dfa0(接收 SparseDenseMatmulFdoConfig*) |
| 缓存默认值 | GetTpuCompEnvWithDefaultValues @ 0x1d73f100(NoDestructor,cxa_guard) |
| Flag→TCE 桥 | OverrideTpuCompEnvByCmdLineFlags @ 0x1d73e640 · SetFieldFromFlagString @ 0x1d73fcc0 |
| Module 获取 | GetTpuCompEnv(HloModule&) @ 0x1d73dd00 → CompilationEnvironments::GetMutableEnv<TCE> @ 0xe6de1e0 |
| Accessor 区段 | 约 130 个 ObjectView<TCE> accessor,位于 0x1d6b0080 .. 0x1d6bf9e0 |
| 置信度 | 除非某行或 callout 另有说明,均为 CONFIRMED(byte-anchored vs decompile) |
1. Proto 身份和注册 {#1-proto-identity-and-registration}
目的
TCE 不是独立的全局 config 对象;它是逐 module 的 compilation environments bag 中的一个条目。理解这一点,是正确获取 TPU 旋钮和追逐一个并不存在的 singleton 之间的区别。本节说明该 message 是什么(一个生成的 proto2 message)、它如何加入 CompilationEnvironments registry,以及消费者如何取回它。
CompilationEnvironment 契约
XLA 将 xla::CompilationEnvironments 定义为一个类型化容器:HloModule 拥有一个这样的容器,后端将其私有配置 message 存入其中,并以该 message 的 proto2::Descriptor 为 key。TCE 通过通用 registry API 参与其中:
xla::CompilationEnvironments
├─ RegisterProcessNewEnvFn(Descriptor*, fn) 0x1e63eb20 ── register a creator keyed by descriptor
│ (guarded by process_new_env_fns_mu; stores into process_new_env_fns map)
├─ AddEnvImpl(Descriptor&, unique_ptr<Message>) 0x1e63ee20 ── install a concrete env into a module's bag
├─ GetMutableEnv<TpuCompilationEnvironment> 0xe6de1e0 ── fetch the TCE by C++ type
└─ CreateFromProto(CompilationEnvironmentsProto) 0x1e63e5a0 ── deserialize a whole bag from the wireGetMutableEnv<T> 是把 C++ 类型映射到其已注册 descriptor 并返回 live instance 的模板(若不存在,则通过已注册 creator 懒惰默认构造)。TCE 是多个已注册 env message 之一;该二进制还携带 GpuCompilationEnvironment(0x1fa44060 等)和 TpuHloModuleBackendConfig(GetMutableEnv<…> @ 0x1c865ac0),它们是同一 registry 中的 sibling。TPU 编译器的私有旋钮位于 TCE arm。
从 HloModule 获取
消费者不读取全局对象;它从正在编译的 module 中读取 TCE。入口点是一小组函数:
GetTpuCompEnv(HloModule&) 0x1d73dd00 ── GetMutableEnv<TCE>( module + 0xF28 )
GetTpuCompEnv(HloModule*) 0x1d73dd20
GetTpuCompEnv(HloInstruction&) 0x1d73dda0 ── via instruction → parent module
GetTpuCompEnv(HloInstruction*) 0x1d73de80
GetTpuCompEnvForAutotuner(...) 0x1305f820 ── autotuner variantGetTpuCompEnv(HloModule&) @ 0x1d73dd00 的反编译是一个单 tail call:
function GetTpuCompEnv(HloModule& m): // 0x1d73dd00
// module field 485 (== module+0xF28) is the CompilationEnvironments*
return CompilationEnvironments::GetMutableEnv<TpuCompilationEnvironment>(m[485])注意 —
+0xF28处的 module 成员(*((QWORD**)module + 485))是该 module 的CompilationEnvironments*。重新实现者必须把 env bag 穿过 HloModule,而不是在 compiler init 时 snapshot 一个扁平 config。每个需要旋钮的 TPU pass 都调用GetTpuCompEnv(module)并读取字段;配置是 module-scoped,而不是 process-scoped。
通过 reflection 的字段访问
flag-binding 路径中的 1121 个字段不是通过 1121 个手写 getter 访问的。TpuCompEnvReflection 层(GetFlagForField @ 0x1d74ad40、GetFieldValue @ 0x1d7523a0、SetEnvField @ 0x1d752ae0、ReadFlag @ 0x1d74af60、ParseFlagFromString @ 0x1d74e8a0)在 descriptor 中按名称/编号查找字段,并泛型地读取或写入。这种 reflection 中介使一个桥能够服务全部 1121 个字段;registry-mediated 路径详见 registry-mediated-flags.md。相比之下,consumer-side hot path 使用约 130 个专用 ObjectView<TpuCompilationEnvironment> accessor(区段 0x1d6b0080 .. 0x1d6bf9e0),它们在固定 struct 偏移处读取字段;这些是 §6 中的 AUTO-resolving accessor。
2. 构造路径
目的
有两个相关的构造函数,它们回答两个不同问题:"给我构建一个新的默认 env"(CreateDefaultTpuCompEnv)与 "给我共享的缓存默认实例"(GetTpuCompEnvWithDefaultValues)。两者都产生一个每个字段都等于其注册 absl::Flag 默认值的 TCE。本节追踪 descriptor-walk-and-copy 机制,以便重新实现者能复现它。
算法 — 默认值实体化
function CreateDefaultTpuCompEnv(SparseDenseMatmulFdoConfig* fdo): // 0x1d73dfa0
env = operator new(0x15E8) // sizeof(TCE) = 5608 B
TpuCompilationEnvironment::ctor(env, /*arena=*/0) // generated proto2 ctor
md = TpuCompilationEnvironment::GetMetadata() // descriptor of the default instance (globals_)
n = md.descriptor.field_count // == 1121
for i in 0 .. n-1: // walk every declared field
fd = md.descriptor.field(i) // proto2::FieldDescriptor*
flag = TpuCompEnvReflection::GetFlagForField(fd) // 0x1d74ad40 — name → FLAGS_<name>
SetFieldFromFlagString(env, fd, flag.default) // 0x1d73fcc0 — copy the flag default union
if fdo != NULL:
// fold the SparseDenseMatmul FDO config into the SparseCore-related fields
ApplyFdoConfig(env, fdo)
return env0x1d73dfa0 的反编译确认了形状:operator new(0x15E8u) → TpuCompilationEnvironment::TpuCompilationEnvironment(env, 0) → 在 TpuCompilationEnvironment_globals_ 默认实例上 GetMetadata() → 一个以 descriptor 字段数量(*(DWORD*)(Metadata + 8))为边界的循环,遍历 proto2::FieldDescriptor 和 absl::CommandLineFlag 对象,并在尾部折入 SparseDenseMatmulFdoConfig* 参数(this/v3)。每字段默认值本身位于 FlagImpl 的默认值 union +0x48 中:那里有 R_X86_64_RELATIVE reloc 表示默认值是一个生成函数(kGenFunc),没有 reloc 则表示 8 字节是内联字面量(kOneWord);所有 1121 个默认值的字节精确统计见 tce-field-offsets-defaults.md。
算法 — 缓存默认值
function GetTpuCompEnvWithDefaultValues(): // 0x1d73f100
static defaults; // NoDestructor<TCE>, cxa_guard'd
if first_call: // __cxa_guard_acquire
tmp = $_0() // 0x1d73f1a0 — build a default TCE on stack
NoDestructor<TCE>::PlacementImpl(&defaults, tmp) // placement-new into static storage
TCE::SharedDtor(tmp) // destroy the temporary
return &defaultslambda $_0 @ 0x1d73f1a0 是 descriptor-walk-and-copy 函数体(与 CreateDefaultTpuCompEnv 相同机制,但没有 FDO 参数)。结果是一个通过引用返回的单一不可变全默认 TCE。这是实际编译的 baseline,之后 OverrideTpuCompEnvByCmdLineFlags @ 0x1d73e640 会应用解析后的 LIBTPU_INIT_ARGS 覆盖(通过 SetFieldFromFlagString @ 0x1d73fcc0 和 OverwriteFieldIfNotDefault @ 0x1d73f360),再之后按版本的 MSA overlay 会重写一组字段。
怪癖 — 默认值是通过对 flag set 做 reflection构建的,而不是手写字段 initializer list。
CreateDefaultTpuCompEnv构造一个空 proto,然后把每个字段注册的absl::Flag默认值复制进去。这就是为什么 TCE schema 和 flag roster 保证 1:1:构造循环若找不到某字段的 flag 就会失败。重新实现者必须为每个字段注册一个 flag(或提供等价的默认值表);这些字段除 flag union 外没有嵌入的 proto-default。(CONFIRMED —0x1d73dfa0中的GetFlagForField循环。)
完整默认值 pipeline 位置
GetTpuCompEnvWithDefaultValues 0x1d73f100 ── all-flag-default baseline (cached)
⊕ OverrideTpuCompEnvByCmdLineFlags 0x1d73e640 ── apply parsed LIBTPU_INIT_ARGS overrides
SetFieldFromFlagString 0x1d73fcc0
OverwriteFieldIfNotDefault 0x1d73f360
⊕ per-TpuVersion MSA overlay (v0/1→jf, v2→cmem, v3→vf, v4/5→gf) ── memory subsystem pages
= effective TCE for this compilation3. 字段组分类法
目的
1121 个字段太多,无法按编号浏览。它们聚集成少量子系统区域;知道某个旋钮属于哪个 cluster,才知道重新实现者如何定位它。本节按名称前缀和功能区域给出分类法。它是地图,不是字典;逐字段名称见 tce-field-dictionary-a.md/tce-field-dictionary-b.md。
Proto 类型统计
1121 个字段按 proto 类型分解如下。每个基础类型在 parse table 中都映射到一个精确的 16 位 FieldEntry.type_card,因此该直方图可从 _table_ @ 0x21cfa9e0 字节精确地重新推导:
| Proto type | Count | type_card | 备注 |
|---|---|---|---|
| bool | 418 | 0x0011 | 153 个默认 true,265 个默认 false |
| message | 349 | 0x0416 | 类型化 sub-message + AutoProto 字段(§4) |
| int64 | 148 | 0x10d1 | thresholds、fuel、ring limits |
| enum | 74 | 0x1891 | 67 个 TristateProto.Value + 7 个 wrapper enum |
| string | 37 | 0x0c15 | 30 个空,7 个非空默认值 |
| float | 34 | 0x1893 | MSA ratios、margins |
| int32 | 32 | 0x1091 | trip counts、levels |
| double | 14 | 0x18d3 | MSA scaling factors |
| uint32 | 11 | 0x0891 | |
| uint64 | 4 | 0x08d1 |
注意 — bool(418)+ message(349)占 schema 的 68%。TCE 主要是一个 boolean feature gate 和类型化 sub-message/AutoProto 旋钮的集合;数值 tunable 是少数。349 个 message-typed 字段在结构上最有意思,其中大多数是通过
AutoOr<T>(§6)解析的AutoProtooneof。完整 type↔type_card 映射和逐类型数量在 tce-field-offsets-defaults.md 上有字节锚定。
按名称前缀 family
TCE 是多个 flag-name family 的落点。两个 proto(TCE vs DebugOptions)之间的拆分由 overview.md 和 flag-families.md 负责;落入 TCE 的 family 是:
| 前缀 family | 子系统 | 详情 |
|---|---|---|
xla_tpu_* | TPU-private core:fusion、scheduling、MSA、collectives、SDC | flag-families.md |
xla_jf_* | Jellyfish backend:LLO codegen、bundle packing、rematerialization | flag-families.md |
xla_sc_* / barna_core_* | SparseCore / BarnaCore embedding | flag-families.md |
xla_msa_* / xla_gf_* / xla_ior_* | Memory-space assignment ratios、register file、CMEM | flag-families.md |
xla_* 的一个子集 | scheduler、profiler、trace toggles,名称与 DebugOptions 共享但存储在 TCE 中 | flag-families.md |
按功能区域
将前缀映射到 prompt 中重新实现者关心的子系统区域,并附带字节锚定的代表字段(field# → name、offset、default,来自 parse table):
| 区域 | 代表字段 | 它们 gate 的内容 |
|---|---|---|
| Scheduler | #31 xla_memory_scheduler(MemoryScheduler enum,默认 DEFAULT),#41 xla_hlo_scheduling_brkga_generation_limit=1200,#151 xla_tpu_licm_analysis_allowance=100000 | memory-pressure scheduler selection、BRKGA genetic-scheduler budget、LICM allowance |
| Fusion | #63 xla_jf_enable_multi_output_fusion=true,#180 xla_tpu_small_operand_count_for_loop_fusion=13,#181 xla_jf_fusion_max_instruction_count_for_window_config=1000,#393 nested-dot-fusion custom ops | multi-output / loop / window fusion eligibility and limits |
| MSA / memory | #280/284/285 jf overlap max/min/pref(32.0/1.0/2.0),#592 xla_tpu_msa_inefficient_use_to_copy_ratio=0.5,#14 xla_tpu_max_cmem_used_by_memory_space_assignment=-1 | memory-space-assignment overlap ratios、copy heuristics、CMEM budget |
| Layout | #758 move_dot_parameters_to_rhs(Tristate ENABLED),#766 enable_large_2nd_minor_layout_for_x8(ENABLED) | dot operand placement、minor-dimension layout |
| SparseCore | #802 enable_offloading_scatter_to_sparsecore(ENABLED),#822/839/860 SC collective-offload(all-gather/all-reduce/reduce-scatter ENABLED),#827 xla_sc_async_wrapper_fusion_type=SINGLE_TPU_CUSTOM_CALL | scatter offload、SC collective offloading、async wrapper fusion |
| Collectives | #55/56/57 arf/ars/agf combiner thresholds=125,829,120 B(120 MiB),#58 xla_jf_crs_combiner_threshold_count=256,#12/13 net-router ring limits=8/16,#149 max_concurrent_send_recv=INT32_MAX | combiner thresholds、net-router rings、send/recv concurrency |
| BarnaCore / embedding | #30 xla_tpu_embedding_table_oblongness_threshold=50.0,#171 internal_embedding_emitter_fraction_vmem_available=0.9 | embedding emitter VMEM budgeting |
| Debug / trace | #32–#39 trace toggles(xla_enable_profiler=true,xla_enable_hlo_trace=true,xla_enable_mxu_trace=false,…),#583 xla_tpu_sdc_checker_checksum_algo,#723 xla_tpu_precision_tracer_mode | profiler/trace emission、SDC checker、precision tracer |
陷阱 — absl flag 默认值是字节权威;help/error 文本不是。两个 flag 的 help string 读作
=false,但它们的FlagImpl+0x48内联字面量是01 00 00 00=true(xla_tpu_rwb_fusion、xla_tpu_accumulate_into_mrb)。重建默认值时,应从 flag/proto initializer(+0x48union)读取,绝不要从 help string 读取。(CONFIRMED — 已在 tce-field-offsets-defaults.md 上按字节修正。)
4. Sub-Message 和 AutoProto 地图 {#4-sub-message-and-autoproto-map}
目的
在 1121 个字段中,349 个是 message-typed。几乎全部不是手写 sub-message,而是 AutoProto oneof,即一个通用 wrapper,保存 30 种可能类型之一的三态值。少数是类型化 helper message(RangeSpecProto)。本节映射 message-typed surface,让重新实现者知道哪些字段是 AutoProto,哪些是类型化 sub-message,以及 AutoProto 的 12 个 message arm 是什么。
类型化 helper message — RangeSpecProto
一小组字段是类型化 RangeSpecProto(它自己的生成 message class,parse table _table_ @ 0x21cf9c90),用于命名 instruction/buffer range 而不是标量的字段,例如 #50 xla_jf_naive_bundle_packer(+0x228)、#60 xla_jf_bounds_check_annotate_only(+0x230)、#65 xla_jf_lsra_v2_alloc_only(+0x238)。它们的默认值是空 message(未设置 range = 按消费者语义 apply-to-all / apply-to-none)。这些字段在 message 区域中位于连续 struct 偏移,因为作为 message 字段,它们以内联方式存储(而不是使用 bool/int 字段那种 8 字节 scalar union)。
AutoProto oneof — 30 个 arm
xla.jellyfish.AutoProto 是一个 30-arm oneof,已从它自己的 parse table AutoProto::_table_ @ 0x21cfa788 逐字节确认。它是大部分 message-typed TCE 字段(约 330 个)的 wrapper:声明为 AutoProto 的字段可以保存 AUTO(unset,oneof-case 0),或者 30 个类型化 alternative 之一。30 个 arm 包括 8 个硬编码 scalar arm(bool/int64/uint64/int32/uint32/double/float/string)、10 个 enum arm,以及12 个 message arm,这些 message arm 通过 parse table 的 12 个 aux pointer 解析到具体 sub-message 类型:
| Arm | Message type | Sub-message _table_ | 管辖内容 |
|---|---|---|---|
| 9 | RepeatedStrings | 0x21cf9d18 | list-valued string knobs |
| 10 | BufferContentsSanitizerConfig | 0x21cf9f58 | buffer sanitizer |
| 13 | CostModelFlagOptions | 0x21cfa170 | cost-model tuning |
| 14 | SparseCoreOffloadingOptions | 0x21cfa110 | SparseCore offload config |
| 15 | ShardyOptions | 0x21cfa260 | Shardy SPMD partitioner |
| 18 | IlpLatencyHidingSchedulerOptions | 0x21cfa308 | ILP latency-hiding scheduler |
| 22 | RepeatedIntegers | 0x21cf9da0 | list-valued integer knobs |
| 23 | EmitterLearnedCostModelOptions | 0x21cff9a8 | learned cost model |
| 26 | AccumulatorTransformations | 0x21cf9c30 | accumulator transform set |
| 27 | SparseCoreAssertLevel | 0x21cfa550 | SC assertion level |
| 28 | BundleInstrumentationOptions | 0x21cfa5b0 | bundle instrumentation |
| 29 | TpuCustomCallMemorySpaceSpec | 0x21cfa708 | custom-call memory-space policy |
30-arm oneof、它的 scalar/enum arm,以及逐 arm message-default 解码由 autoproto-message-arms.md 负责。AutoProto 字段的默认值是空 AutoProto(oneof-case 0 = AUTO/unset),因此值由消费者的解析极性提供(§6)。
嵌套 policy sub-message — MsaReservationPolicy、HbmPolicy
memory subsystem 其他地方命名的 MSA reservation 和 HBM policy 嵌套在 TpuCustomCallMemorySpaceSpec message arm 下(arm 29),而不是顶层 TCE 字段。二进制将它们携带为 TpuCustomCallMemorySpaceSpec_MsaReservationPolicy(Clear @ 0x1db25da0,ctor @ 0x1db25c80)和 TpuCustomCallMemorySpaceSpec_HbmPolicy(ctor @ 0x1db25ea0,GetClassData @ 0x1db25ee0),并由 TpuCustomCallMemorySpacePolicy::RunMsaReservationPolicy @ 0x110367c0 和 …::RunHbmPolicy @ 0x11038120 消费。
怪癖 —
MsaReservationPolicy和HbmPolicy是 AutoProto message arm(TpuCustomCallMemorySpaceSpec)的嵌套 sub-message,位于 TCE 下两层。重新实现者无法通过顶层 TCE 字段号到达它们;只有当相关 AutoProto 字段的TpuCustomCallMemorySpaceSpecarm 被设置时,它们才会被填充,并驱动RunMsaReservationPolicy/RunHbmPolicy在 MSA 期间应用的 custom-call memory-space assignment。(CONFIRMED — mangledTpuCustomCallMemorySpaceSpec_MsaReservationPolicy/_HbmPolicysymbols。)
5. Wire Format、Presence 和内存布局
目的
TCE 是一个真正的 proto2 message:它会序列化、携带 has-bit presence,并通过 module 的 CompilationEnvironments bag 往返到 wire。重新实现者需要布局常量来读取序列化 env,并判断哪些字段被用户触碰,哪些保留默认值。本节给出 parse-table 头部和 serialization 入口,两者均有字节锚定。
Parse-table 头部 — 布局 oracle
TpuCompilationEnvironment::_table_ @ 0x21cfa9e0 是驱动 fast-path parsing 的 TcParseTableBase,其头部常量也充当内存布局地图:
TcParseTableBase @ 0x21cfa9e0
+0x00 has_bits_offset = 16 (0x10) ── presence bitfield starts at struct +0x10
+0x02 extension_offset = 0 ── no proto extensions
+0x04 max_field_number = 1218 (0x4c2) ── highest declared field# (gaps tombstoned)
+0x08 fast_idx_mask = 0xf8
+0x10 field_entries_offset = 0x370 ── FieldEntry[1121] starts here
+0x14 num_field_entries = 1121
+0x16 num_aux_entries = 349 ── == the 349 message-typed fields
+0x18 aux_offset = 0x3800
sizeof(TCE) = 0x15e8 (5608 B); field data region +0xA8 .. +0x15E0每个 FieldEntry 为 12 字节 {uint32 offset; uint32 has_idx; uint16 aux_idx; uint16 type_card},数组按字段号升序排序,因此 entry[i] 是第 i 个存活字段,且 struct_offset(field N) = FieldEntry[index_of(N)].offset。field#→offset 机制和确定性偏移公式由 tce-field-offsets-defaults.md 负责;本页只定位该头部。
Presence 模型
has_bits_offset = 16 意味着每个字段都有一个显式 presence bit,打包在 +0x10 .. ~+0xB0 区域中(每字段一个,按 FieldEntry.has_idx 顺序)。这使 override 路径能区分“用户把这个旋钮设为 false”和“保留默认 false”:OverwriteFieldIfNotDefault @ 0x1d73f360 在决定按版本 overlay 是否可以重写字段之前,会读取 has-bit。丢弃 proto2 presence 的重新实现者无法正确实现 ⊕-overlay 逻辑。
注意 — 精确的
has_idx → (struct word, bit)打包没有从二进制重新走查;FieldEntry.has_idx值已恢复,布局遵循 protobuf 从has_bits_offset = 16开始覆盖+0x10 .. ~+0xB0区域(约 1121 bits)的标准顺序打包。presence-region 边界和逐字段has_idx已 CONFIRMED;精确 bit assignment 为 HIGH-confidence(标准打包),但未逐字节验证。(HIGH)
序列化往返
TCE 在 wire 上作为 CompilationEnvironmentsProto 的一个条目传输。CompilationEnvironments::CreateFromProto @ 0x1e63e5a0 反序列化整个 bag:对每个序列化 env,它通过 message descriptor 查找已注册 creator(来自 §1 的 RegisterProcessNewEnvFn map),并重建类型化 message。因此,只要双方都注册了 TCE descriptor,一个 host(例如 autotuner run)生成的 TCE 就可以序列化进 module 的 CompilationEnvironmentsProto,并在另一个 host 上忠实重建。
6. AutoOr<T> Wrapper {#6-the-autoort-wrapper}
目的
约 330 个 AutoProto 字段不携带普通值,而是携带三态(布尔的 AUTO / ENABLED / DISABLED,或其他类型的 AUTO-vs-explicit),具体编译值通过 AutoOr<T> accessor 在消费者处懒惰计算,其极性就是文档化默认值。本节定位 wrapper;字节精确 resolver 和逐旋钮极性统计由 autoproto-autoor-resolution.md 负责。
它如何应用到 TCE 字段
消费者通过 ObjectView<TpuCompilationEnvironment> accessor(区段 0x1d6b0080 .. 0x1d6bf9e0)读取 AutoProto 字段。accessor 在字段的 struct 偏移处取出 AutoProto,检查 oneof discriminator,并应用逐旋钮极性:
function read_auto_bool_field(ObjectView<TCE> env, offset OFF): // one of the 0x1d6b* accessors
p = env.field_at(OFF) // AutoProto sub-message
r = AutoOr<bool>::FromProtoOrDie(p) // 0xf795300 — (present<<8)|val, AUTO ⇒ 0
if KNOB_POLARITY == AUTO_off: // only explicit true enables
return (r == (1<<8 | 1)) // present AND true
else: # AUTO_on // only explicit false disables
return (r != (1<<8 | 0)) // not (present AND false)因此,同样的全 AUTO 存储会根据旋钮作者的 idiom 产生相反默认值。具有 AUTO=off 极性的字段默认 OFF;具有 AUTO=on 极性的字段默认 ON,两者都由同一个空 AutoProto 支撑。还存在第二种存储类别:67 个字段是内联 TristateProto.Value enum(由 cmpl $0x2, +OFF 读取的 4 字节字段,其中 ENABLED==2),不是 AutoProto wrapper。完整极性统计(offset → polarity)、int64-sentinel idiom 和逐类型打包见 autoproto-autoor-resolution.md。
Wrapper enum
除 67 个 TristateProto.Value 字段外,7 个 TCE enum 字段使用专用 wrapper enum,其默认值为非零值。例如 #631 xla_tpu_register_selection_policy = 6(DISREGARD_RECENTLY_USED)、#827 xla_sc_async_wrapper_fusion_type = 3(SINGLE_TPU_CUSTOM_CALL)、#132 xla_tpu_verify_or_assign_tiling_before_lowering = 1(VERIFY)。把解析后的 flag 值映射到正确 TCE oneof arm 的桥,是 NormalizeFieldType<T> / AutoOr<T> 模板 family(每种 wrapper type 一个实例)。完整的 17 个 wrapper-enum 值表见 autoproto-autoor-resolution.md。
陷阱 — 不存在 eager 的“resolve all knobs” pass。
AutoOr<T>解析在每次 accessor 调用时运行,并在字段未设置时回退到空 AutoProto 默认值。若重新实现者在 init 时实体化一个扁平bool/intconfig struct 并在之后读取它,会悄悄地为约 330 个 AutoProto 字段提供错误默认值(并错过按版本 overlay)。配置是按消费者、按字段懒惰解析的。(CONFIRMED —FromProtoOrDie @ 0xf795300packing。)
相关组件
| 组件 | 关系 |
|---|---|
xla::DebugOptions(debugoptions-proto.md) | 后端共享的 sibling proto;携带通用/dump/GPU/CPU 字段,在 TPU 上大多惰性 |
CompilationEnvironments::GetMutableEnv<TCE> @ 0xe6de1e0 | 按 C++ 类型从 module 的 env bag 获取 TCE |
CreateDefaultTpuCompEnv @ 0x1d73dfa0 | 通过 descriptor-walk + flag-default copy 构建默认 TCE |
GetTpuCompEnvWithDefaultValues @ 0x1d73f100 | once-guarded 的缓存全默认 TCE 实例 |
OverrideTpuCompEnvByCmdLineFlags @ 0x1d73e640 | 在默认值之上应用解析后的 LIBTPU_INIT_ARGS overrides |
AutoProto::_table_ @ 0x21cfa788 | 约 330 个 AutoProto 字段背后的 30-arm oneof parse table |
TpuHloModuleBackendConfig / GpuCompilationEnvironment | 同一 CompilationEnvironments registry 中的 sibling registered env message |
交叉引用
- overview.md — 三层 flag→proto→effective-value pipeline;TCE 是 Stage 2b
- debugoptions-proto.md — 后端共享的
xla::DebugOptionsschema;TCE 的通用 sibling - tce-field-dictionary-a.md / tce-field-dictionary-b.md — 逐字段 name/type/meaning 字典(完整 1121 字段列表)
- tce-field-offsets-defaults.md — 字节精确的 field#→struct-offset→default 参考和 parse-table 机制
- autoproto-autoor-resolution.md —
AutoOr<T>三态 resolver、约 130 个 accessor 的极性统计、int64 sentinels 和 17 个 wrapper-enum 表 - autoproto-message-arms.md — 30-arm AutoProto oneof 及其 12 个 message arm 的细节
- registry-mediated-flags.md —
TpuCompEnvReflectionflag↔field 桥 - flag-families.md — 哪些 name-prefix family 落入 TCE、DebugOptions 或 standalone flags
- xla-flag-atlas.md — 与 TCE 字段集 1:1 镜像的约 2107 个名称的 flag surface