Skip to content

TpuCompilationEnvironment

本页所有符号、地址和 struct 偏移均适用于 libtpu-0.0.40-cp314 wheel 中的 libtpu.so(构建 libtpu_lts_20260413_b_RC00,build-id md5 89edbbe81c5b328a958fe628a9f2207d,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::Flag default union 来构建默认值,以及缓存默认值如何只实体化一次。
  • 字段组分类法 — 1121 个字段聚集到哪些子系统区域(让重新实现者知道某个旋钮位于哪里),以及界定 schema 的 proto 类型统计。
  • sub-message + AutoProto 地图 — 类型化 sub-message(RangeSpecProto、12 个 AutoProto message arm、嵌套 TpuCustomCallMemorySpaceSpec policy)以及 AutoOr<T> 三态适用的位置。
Messagexla::jellyfish::TpuCompilationEnvironment(package xla.jellyfish,proto2)
字段数量1121 个存活字段(最大 field# 1218,gap 已 tombstoned)
sizeof0x15E8 (5608 B);字段数据区域 +0xA8 .. +0x15E0
Parse table_table_ @ 0x21cfa9e0TcParseTableBase + 1121-entry FieldEntry 数组 @ +0x370
默认实例TpuCompilationEnvironment_globals_ @ 0x227b87e0
Class dataTpuCompilationEnvironment_class_data_ @ 0x223c96a0sizeof 0x15e8 at +0x20
构造默认值CreateDefaultTpuCompEnv @ 0x1d73dfa0(接收 SparseDenseMatmulFdoConfig*
缓存默认值GetTpuCompEnvWithDefaultValues @ 0x1d73f100NoDestructorcxa_guard
Flag→TCE 桥OverrideTpuCompEnvByCmdLineFlags @ 0x1d73e640 · SetFieldFromFlagString @ 0x1d73fcc0
Module 获取GetTpuCompEnv(HloModule&) @ 0x1d73dd00CompilationEnvironments::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 参与其中:

text
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 wire

GetMutableEnv<T> 是把 C++ 类型映射到其已注册 descriptor 并返回 live instance 的模板(若不存在,则通过已注册 creator 懒惰默认构造)。TCE 是多个已注册 env message 之一;该二进制还携带 GpuCompilationEnvironment0x1fa44060 等)和 TpuHloModuleBackendConfigGetMutableEnv<…> @ 0x1c865ac0),它们是同一 registry 中的 sibling。TPU 编译器的私有旋钮位于 TCE arm。

从 HloModule 获取

消费者不读取全局对象;它从正在编译的 module 中读取 TCE。入口点是一小组函数:

text
GetTpuCompEnv(HloModule&)        0x1d73dd00  ── GetMutableEnv<TCE>( module + 0xF28 )
GetTpuCompEnv(HloModule*)        0x1d73dd20
GetTpuCompEnv(HloInstruction&)   0x1d73dda0  ── via instruction → parent module
GetTpuCompEnv(HloInstruction*)   0x1d73de80
GetTpuCompEnvForAutotuner(...)   0x1305f820  ── autotuner variant

GetTpuCompEnv(HloModule&) @ 0x1d73dd00 的反编译是一个单 tail call:

c
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 @ 0x1d74ad40GetFieldValue @ 0x1d7523a0SetEnvField @ 0x1d752ae0ReadFlag @ 0x1d74af60ParseFlagFromString @ 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 机制,以便重新实现者能复现它。

算法 — 默认值实体化

c
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 env

0x1d73dfa0 的反编译确认了形状:operator new(0x15E8u)TpuCompilationEnvironment::TpuCompilationEnvironment(env, 0) → 在 TpuCompilationEnvironment_globals_ 默认实例上 GetMetadata() → 一个以 descriptor 字段数量(*(DWORD*)(Metadata + 8))为边界的循环,遍历 proto2::FieldDescriptorabsl::CommandLineFlag 对象,并在尾部折入 SparseDenseMatmulFdoConfig* 参数(this/v3)。每字段默认本身位于 FlagImpl 的默认值 union +0x48 中:那里有 R_X86_64_RELATIVE reloc 表示默认值是一个生成函数(kGenFunc),没有 reloc 则表示 8 字节是内联字面量(kOneWord);所有 1121 个默认值的字节精确统计见 tce-field-offsets-defaults.md

算法 — 缓存默认值

c
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 &defaults

lambda $_0 @ 0x1d73f1a0 是 descriptor-walk-and-copy 函数体(与 CreateDefaultTpuCompEnv 相同机制,但没有 FDO 参数)。结果是一个通过引用返回的单一不可变全默认 TCE。这是实际编译的 baseline,之后 OverrideTpuCompEnvByCmdLineFlags @ 0x1d73e640 会应用解析后的 LIBTPU_INIT_ARGS 覆盖(通过 SetFieldFromFlagString @ 0x1d73fcc0OverwriteFieldIfNotDefault @ 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 位置

text
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 compilation

3. 字段组分类法

目的

1121 个字段太多,无法按编号浏览。它们聚集成少量子系统区域;知道某个旋钮属于哪个 cluster,才知道重新实现者如何定位它。本节按名称前缀和功能区域给出分类法。它是地图,不是字典;逐字段名称见 tce-field-dictionary-a.md/tce-field-dictionary-b.md

Proto 类型统计

1121 个字段按 proto 类型分解如下。每个基础类型在 parse table 中都映射到一个精确的 16 位 FieldEntry.type_card,因此该直方图可从 _table_ @ 0x21cfa9e0 字节精确地重新推导:

Proto typeCounttype_card备注
bool4180x0011153 个默认 true,265 个默认 false
message3490x0416类型化 sub-message + AutoProto 字段(§4
int641480x10d1thresholds、fuel、ring limits
enum740x189167 个 TristateProto.Value + 7 个 wrapper enum
string370x0c1530 个空,7 个非空默认值
float340x1893MSA ratios、margins
int32320x1091trip counts、levels
double140x18d3MSA scaling factors
uint32110x0891
uint6440x08d1

注意 — bool(418)+ message(349)占 schema 的 68%。TCE 主要是一个 boolean feature gate 和类型化 sub-message/AutoProto 旋钮的集合;数值 tunable 是少数。349 个 message-typed 字段在结构上最有意思,其中大多数是通过 AutoOr<T>§6)解析的 AutoProto oneof。完整 type↔type_card 映射和逐类型数量在 tce-field-offsets-defaults.md 上有字节锚定。

按名称前缀 family

TCE 是多个 flag-name family 的落点。两个 proto(TCE vs DebugOptions)之间的拆分由 overview.mdflag-families.md 负责;落入 TCE 的 family 是:

前缀 family子系统详情
xla_tpu_*TPU-private core:fusion、scheduling、MSA、collectives、SDCflag-families.md
xla_jf_*Jellyfish backend:LLO codegen、bundle packing、rematerializationflag-families.md
xla_sc_* / barna_core_*SparseCore / BarnaCore embeddingflag-families.md
xla_msa_* / xla_gf_* / xla_ior_*Memory-space assignment ratios、register file、CMEMflag-families.md
xla_* 的一个子集scheduler、profiler、trace toggles,名称与 DebugOptions 共享但存储在 TCE 中flag-families.md

按功能区域

将前缀映射到 prompt 中重新实现者关心的子系统区域,并附带字节锚定的代表字段(field# → name、offset、default,来自 parse table):

区域代表字段它们 gate 的内容
Scheduler#31 xla_memory_schedulerMemoryScheduler enum,默认 DEFAULT),#41 xla_hlo_scheduling_brkga_generation_limit=1200,#151 xla_tpu_licm_analysis_allowance=100000memory-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 opsmulti-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=-1memory-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_CALLscatter 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_MAXcombiner 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.9embedding 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_modeprofiler/trace emission、SDC checker、precision tracer

陷阱 — absl flag 默认值是字节权威;help/error 文本不是。两个 flag 的 help string 读作 =false,但它们的 FlagImpl+0x48 内联字面量是 01 00 00 00 = truexla_tpu_rwb_fusionxla_tpu_accumulate_into_mrb)。重建默认值时,应从 flag/proto initializer(+0x48 union)读取,绝不要从 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 类型:

ArmMessage typeSub-message _table_管辖内容
9RepeatedStrings0x21cf9d18list-valued string knobs
10BufferContentsSanitizerConfig0x21cf9f58buffer sanitizer
13CostModelFlagOptions0x21cfa170cost-model tuning
14SparseCoreOffloadingOptions0x21cfa110SparseCore offload config
15ShardyOptions0x21cfa260Shardy SPMD partitioner
18IlpLatencyHidingSchedulerOptions0x21cfa308ILP latency-hiding scheduler
22RepeatedIntegers0x21cf9da0list-valued integer knobs
23EmitterLearnedCostModelOptions0x21cff9a8learned cost model
26AccumulatorTransformations0x21cf9c30accumulator transform set
27SparseCoreAssertLevel0x21cfa550SC assertion level
28BundleInstrumentationOptions0x21cfa5b0bundle instrumentation
29TpuCustomCallMemorySpaceSpec0x21cfa708custom-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 — MsaReservationPolicyHbmPolicy

memory subsystem 其他地方命名的 MSA reservation 和 HBM policy 嵌套在 TpuCustomCallMemorySpaceSpec message arm 下(arm 29),而不是顶层 TCE 字段。二进制将它们携带为 TpuCustomCallMemorySpaceSpec_MsaReservationPolicyClear @ 0x1db25da0ctor @ 0x1db25c80)和 TpuCustomCallMemorySpaceSpec_HbmPolicyctor @ 0x1db25ea0GetClassData @ 0x1db25ee0),并由 TpuCustomCallMemorySpacePolicy::RunMsaReservationPolicy @ 0x110367c0…::RunHbmPolicy @ 0x11038120 消费。

怪癖 — MsaReservationPolicyHbmPolicy 是 AutoProto message armTpuCustomCallMemorySpaceSpec)的嵌套 sub-message,位于 TCE 下两层。重新实现者无法通过顶层 TCE 字段号到达它们;只有当相关 AutoProto 字段的 TpuCustomCallMemorySpaceSpec arm 被设置时,它们才会被填充,并驱动 RunMsaReservationPolicy/RunHbmPolicy 在 MSA 期间应用的 custom-call memory-space assignment。(CONFIRMED — mangled TpuCustomCallMemorySpaceSpec_MsaReservationPolicy / _HbmPolicy symbols。)


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,其头部常量也充当内存布局地图:

text
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(来自 §1RegisterProcessNewEnvFn 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,并应用逐旋钮极性:

c
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 = 6DISREGARD_RECENTLY_USED)、#827 xla_sc_async_wrapper_fusion_type = 3SINGLE_TPU_CUSTOM_CALL)、#132 xla_tpu_verify_or_assign_tiling_before_lowering = 1VERIFY)。把解析后的 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/int config struct 并在之后读取它,会悄悄地为约 330 个 AutoProto 字段提供错误默认值(并错过按版本 overlay)。配置是按消费者、按字段懒惰解析的。(CONFIRMED — FromProtoOrDie @ 0xf795300 packing。)


相关组件

组件关系
xla::DebugOptionsdebugoptions-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 @ 0x1d73f100once-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

交叉引用