MSA 每版本默认值

本页中的所有地址、偏移和 flag 值均适用于 libtpu 0.0.40（libtpu-0.0.40-cp314）、libtpu.so build-id 89edbbe81c5b328a958fe628a9f2207d。其他 wheel，以及任何面向 v5e/v6e 的构建，可能不同。

摘要

Memory-Space Assignment（MSA）是 TPU 编译器第 7 阶段的 alternate-memory 放置 pass（见 msa-overview.md），它并不是由单个硬编码 option struct 驱动。它的调优旋钮（overlap-to-async-copy ratios、outstanding-copy caps、repack/retry counts、inefficient-use skip threshold、cross-program-prefetch enables，以及 alternate-memory byte budget）会按 TPU 版本从分层的 absl 命令行 flags 解析。每个 tpu_version（tpu::TpuVersion enum 值：0=Jellyfish，1=Dragonfish，2=Pufferfish，3=Viperfish，4=Ghostlite，5=6acc60406）被映射到四个 flag-prefix families 之一：jf（覆盖 Jellyfish 0 + Dragonfish 1）、cmem（CMEM 层级，Pufferfish 2）、vf（Viperfish 3）、gf（覆盖 Ghostlite 4 + 6acc60406 5）；另有最后应用的全局 xla_msa_* fallback。因此任意旋钮的有效值都是 (version → family selection) × (family flag default) 的笛卡尔积。

重新实现者的自然假设，也就是生成本页的任务中提出的假设，是每版本数字存在于随 chip-config 资源一起发布的 *_compilation_environment.binarypb overlay 中。并非如此。 这个构建中不存在这样的 overlay 资源；filewrapper 目录表和 .rodata 只携带 chip_parts、chip_configs、bootloaders、accuracy tables、route caches、ICI-resiliency、CSR/error blobs 和一个 timezone blob。字面默认值实际位于 .data 中的 absl::flags_internal::Flag<T> 对象内，运行时在首次使用时将它们物化到默认 TpuCompilationEnvironment proto 中。由于此构建里的 memory_space_assignment.proto 和 tpu_compilation_environment.proto 都是 proto3，两个 descriptor 都不存储 default_value 字符串，这也正是 editions/text-format printer 路径不显示默认值的原因。

本页负责记录字节精确的每 family 数值表和两阶段解析路径：（1）GetTpuCompEnvWithDefaultValues() 遍历 descriptor，并把每个 absl flag 当前默认值复制到匹配字段；（2）ComputeMemorySpaceAssignmentOptions() 按版本选择 family-specific flag，并以全局 xla_msa_* flag 作为 fallback，通过 OverwriteFieldIfNotDefault 完成。version→family 选择是 IsMemorySpaceAssignmentEnabled 内的 6 项 jump table。

重新实现契约为：

version→family jump table（Target+0x398 处的 tpu_version，cases 0/1→jf，2→cmem，3→vf，4/5→gf），以及 OverwriteFieldIfNotDefault(family_flag, msa_fallback) 解析规则。
字面每 family 默认值：三个 overlap ratios（min/preferred 统一，max split 为 32×-jf 对 8×-rest）、outstanding-copy caps（4-jf 对 40-rest）、repacks=4、retries=2、inefficient-use=0.5、cross-program-prefetch=1。
alternate-memory byte budget 不是 flag：byte-cap flags 都默认 -1（“从 Target 推导”），因此 Options.max_size_in_bytes 来自 chip_parts，不是字面值。
absl::flags_internal::Flag<T> 对象布局（+0x48 default-value union：floats/bools 使用 kGenFunc，int caps 使用 kOneWord inline literal）。


Version→family gate	`xla::jellyfish::IsMemorySpaceAssignmentEnabled` @ `0x12fc1280`
Version jump table	`0xae09ac8` — 6 项（v0/1→jf，v2→cmem，v3→vf，v4/5→gf）
Per-version selector	`xla::jellyfish::ComputeMemorySpaceAssignmentOptions` @ `0x12fc1440`（55 次 `OverwriteFieldIfNotDefault` 调用）
Default materializer	`xla::jellyfish::GetTpuCompEnvWithDefaultValues` @ `0x1d73f100`
Resolution primitive	`xla::jellyfish::OverwriteFieldIfNotDefault(family_flag, fallback, env&)` @ `0x1d73f360`
`tpu_version` field	`Target + 0x398`（`int32`）
Default-value source	`absl::flags_internal::Flag<T>` object `+0x48`（union：`kGenFunc` / `kOneWord`）
Env binarypb overlay	此构建中无（没有 `*_compilation_environment.binarypb`）

代号 Family 与版本门控

目的

MSA 做出的第一个决定是它是否运行，并且通过同一路径决定由哪个 family 的 flags 管辖。IsMemorySpaceAssignmentEnabled 读取 tpu_version，并通过 jump table 跳到 family-specific *_memory_space_assignment enable flag，同时以 xla_msa_enable 作为全局 fallback。随后相同的 family selection（jf / cmem / vf / gf）会管辖 ComputeMemorySpaceAssignmentOptions 中的每个数值旋钮。

Family 前缀

前缀	代号	`tpu_version`	角色
`jf`	Jellyfish + Dragonfish	0, 1	最老 family；唯一具有非默认 ratio/cap 值的 family。Dragonfish (1) 共享 Jellyfish 的 `jf` flags，与其共享的 encoder/bundle-restrictions 相匹配
`cmem`	Pufferfish	2	为 Pufferfish 选择的 CMEM 层级；携带反向的 `min_async_copy_to_overlap_ratio` twin
`vf`	Viperfish	3
`gf`	Ghostlite + 6acc60406	4, 5	两代共享 `gf` flags，与共享的 `gxc` ISA family（Ghostlite 为 `gxc::glc`，6acc60406 为 `gxc::gfc`）相匹配
`msa`	global fallback	—	`OverwriteFieldIfNotDefault` 最后应用的 `xla_msa_*` flags
`zf`	forward stub	—（未门控）	`xla_zf_vmem_max_outstanding_*` 存在，但 `zf` 不在 0..5 表中

QUIRK — zf 是此二进制中的真实 flag 前缀（xla_zf_vmem_max_outstanding_prefetches HelpGen @ 0x1d72ef40，_evictions @ 0x1d72f000，二者默认 40），但它不在 0..5 version gate 中，也没有 overlap-ratio / repack / retry 变体。它是尚未接入任何 tpu_version 的 forward-family stub。重新实现者应保留 flag 名称，但不得把任何当前版本路由到它。

算法

门控对 *(Target + 0x398)（反编译器命名为 a2 + 920 = 0x398）进行 switch。每个 case 都先以全局 xla_msa_enable flag、再以 family-specific *_memory_space_assignment flag 作为每版本 override 调用 OverwriteFieldIfNotDefault；解析结果落入本地 TpuCompilationEnvironment，并在 LABEL_11 读回。

function IsMemorySpaceAssignmentEnabled(out, Target, env_view, HloModule):  // 0x12fc1280
    if IsPassDisabled("memory-space-assignment", env_view):     // 0x12fc1280 entry guard
        out.enabled = false; return                             //   pass explicitly disabled

    TpuCompilationEnvironment local;                            // built from defaults
    switch (*(int32*)(Target + 0x398)):                         // tpu_version
        case 0: case 1:                                         // jf — Jellyfish + Dragonfish
            OverwriteFieldIfNotDefault("xla_msa_enable",
                "xla_jf_vmem_memory_space_assignment", local)
        case 2:                                                 // cmem — Pufferfish
            OverwriteFieldIfNotDefault("xla_msa_enable",
                "xla_tpu_cmem_memory_space_assignment", local)
        case 3:                                                 // vf — Viperfish
            OverwriteFieldIfNotDefault("xla_msa_enable",
                "xla_vf_vmem_memory_space_assignment", local)
        case 4: case 5:                                         // gf — Ghostlite + 6acc60406
            OverwriteFieldIfNotDefault("xla_msa_enable",
                "xla_gf_vmem_memory_space_assignment", local)
        default:                                                // unknown version
            // LABEL_11: fall through to the Tristate read below

    // LABEL_11 — read the resolved enable:
    if local.msa_enable_tristate != UNSET:                      // local field v11
        out.enabled = (local.msa_enable_tristate == ENABLED)    //   Tristate==2
    else:
        out.enabled = (*(int32*)(Target + 0x398) >= 2)          //   default: on for v>=2
```text

> **NOTE —** 解析原语名为 `OverwriteFieldIfNotDefault(global, override, env)`：它先把全局默认值写入字段，然后仅当 family-specific flag 非默认时，才用该 flag 的值覆盖。反编译调用将全局项（`xla_msa_enable`，长度 14）作为参数 1，将 family flag（`xla_jf/vf/gf_vmem_memory_space_assignment` 长度 35，`xla_tpu_cmem_memory_space_assignment` 长度 36）作为 override，逐字节匹配 flag-name `.rodata` 长度。`AddSourceLocationImpl` 错误路径记录的源文件是 `platforms/xla/service/jellyfish/memory_space_assignment_util.cc`。

### enable 默认值

五个 `*_memory_space_assignment` family flags 全部默认 **1（enabled）**。全局 `xla_msa_enable` 是 **Tristate**，其 default-gen（`AbslFlagDefaultGenForxla_msa_enable::Gen` @ `0x1d705500`）写入单字节 **2 = ENABLED**。因此 MSA 对每个被门控版本默认开启；`default:` fallthrough 还会额外为任何未知 `tpu_version >= 2` 启用它。

| Flag | 类型 | 默认 |
|---|---|---|
| `xla_jf_vmem_memory_space_assignment` | bool | **1** |
| `xla_vf_vmem_memory_space_assignment` | bool | **1** |
| `xla_gf_vmem_memory_space_assignment` | bool | **1** |
| `xla_tpu_cmem_memory_space_assignment` | bool | **1** |
| `xla_msa_enable` (global) | Tristate | **2 (ENABLED)** |

---

## 每 Family 数值默认值

这些是重新实现者必须复现的字面值。每个 float 都从其 default-gen body（`mov dword[rdi], imm32; ret`）按字节精确物化；每个 int cap 是 `Flag<T>+0x48` 处的 inline `kOneWord` literal。Confidence 列标记验证方式。

### Overlap-to-async-copy ratios

三个 float 会馈入 `PrefetchIntervalPicker` 构造函数（`@0x1dcd6b60`）：`picker+0x80` ← min，`picker+0x84` ← preferred，`picker+0x88` ← 由 preferred × max 派生的乘积。`min` 为最晚合法 prefetch 时间播种（`use_time − ⌈min × async_copy_elapsed⌉`）；`max` 约束 prefetch 可被提前到使用点之前多远。

| Option | jf | vf | gf | cmem | msa |
|---|---:|---:|---:|---:|---:|
| `min_overlap_to_async_copy_ratio` | 1.0 | 1.0 | 1.0 | 1.0 | 1.0 |
| `preferred_overlap_to_async_copy_ratio` | 2.0 | 2.0 | 2.0 | 2.0 | 2.0 |
| `max_overlap_to_mem_size_async_copy_ratio` | **32.0** | 8.0 | 8.0 | 8.0 | 8.0 |

> **QUIRK —** 只有 **max** ratio 因 family 而异，并且只有 Jellyfish 与众不同：在 `jf` 上，buffer 最多可提前到其自身 async-copy elapsed time 的 **32×** 被 prefetch；之后所有 family 都只有 **8×**。`min`（1.0）和 `preferred`（2.0）在所有 families 上一致。字节证据：`jf` max default-gen @ `0x1d72cae0` 写入 `0x42000000`（=32.0f）；`gf` max @ `0x1d72dfa0` 和 `vf` max @ `0x1d72d4a0` 均写入 `0x41000000`（=8.0f）；`jf` min @ `0x1d72cc80` 写入 `0x3F800000`（=1.0f）；`jf` preferred @ `0x1d72cd60` 写入 `0x40000000`（=2.0f）。`cmem` 额外携带一个反向 twin：`xla_tpu_cmem_min_async_copy_to_overlap_ratio` = 1.0（Gen @ `0x1d72fb20` 写入 `0x3F800000`）。

### Outstanding-copy caps 与 repack/retry counts

outstanding caps 是 `AsynchronousCopyResource` time-bucket 模型强制执行的上限：会超过 cap 的候选 copy 返回 `FailOutOfAsyncCopies`（status `0x10`）。这些值作为 `kOneWord` inline literals 存储（没有 default-gen function），并由 `ComputeMemorySpaceAssignmentOptions` 中引用的 family-paired flag 名称确认。

| Option | jf | vf | gf | zf | cmem | msa |
|---|---:|---:|---:|---:|---:|---:|
| `max_outstanding_prefetches` | **4** | 40 | 40 | 40 | 40 | 40 |
| `max_outstanding_evictions` | **4** | 40 | 40 | 40 | 40 | 40 |
| `max_repacks` | 4 | 4 | 4 | — | 4 | 4 |
| `max_retries` | 2 | 2 | 2 | — | 2 | 2 |

> **QUIRK —** outstanding-copy budget 将 Jellyfish 与其他所有 family 分出 **10×** 差距：`jf` 上并发 prefetches/evictions 为 4，而之后所有 family（以及 `zf`）为 40。`max_repacks`（4）和 `max_retries`（2）在所有 families 上一致。int caps 是 `Flag<T>+0x48` 处的 inline `kOneWord` literals，不是 default-gen functions，因此反编译中不会以 `Gen` 符号出现；它们通过 `ComputeMemorySpaceAssignmentOptions` @ `0x12fc1440` 中 family-paired `OverwriteFieldIfNotDefault` 参数对（`xla_{jf,vf,gf,tpu_cmem,msa}_…_max_outstanding_prefetches`、`…_evictions`、`…_max_repacks`、`…_max_retries`）的名称得到确认。

### Cross-program prefetch

Cross-program prefetch 是除 `max_overlap` 外唯一一个其 *enable* 因 family 而异的旋钮。cap 本身是全局的（1），但 per-family `*_vmem_enable_cross_program_prefetch` bool 在 Jellyfish 上为 **0**，在 vf/gf 上为 1。

| Flag | jf | vf | gf | cmem | msa |
|---|---:|---:|---:|---:|---:|
| `<family>_vmem_enable_cross_program_prefetch` | **0** | 1 | 1 | — | — |
| `<family>_vmem_enable_cross_program_prefetch_freeing` | — | 1 | 1 | — | 1 |
| `<family>_vmem_enable_while_redundant_eviction_elimination` | 1 | 1 | 1 | 1 | 1 |
| `<family>_vmem_default_cross_program_prefetch_heuristic` | 0 | 0 | 0 | — | 0 |

`msa_enable_cross_program_prefetch_freeing` default-gen @ `0x1d705e00` 写入字节 **1**；`jf`/`vf`/`gf` 的 `enable_cross_program_prefetch` 和 `_freeing` flags 存在（HelpGen 符号存在，例如 `jf` enable @ `0x1d72c7e0`、`vf` enable @ `0x1d72d4c0`、`gf` enable @ `0x1d72dfc0`），但使用 `kOneWord`/inline defaults 而不是 `Gen` bodies，因此 per-family 0/1 分裂为 HIGH（不是来自 `Gen` 的字节精确值）。

family flags 之下的全局共享 bools：

| Global flag | 默认 |
|---|---:|
| `xla_tpu_enable_cross_program_prefetch_freeing` | 1 |
| `xla_enable_cross_program_prefetch` | 1 |
| `xla_default_cross_program_prefetch_heuristic` | 0 |

---

## Singleton 与全局默认值

若干 MSA 旋钮没有 per-family 变体，而是由单个全局 flag 管辖所有版本。byte-cap flags 最重要：它们全部默认 **-1**，表示“无 override；从 `Target` 推导 budget”。

### Float 与 integer globals

| Flag | 类型 | 默认 |
|---|---|---:|
| `xla_tpu_msa_inefficient_use_to_copy_ratio` | float | **0.5** |
| `xla_msa_max_cross_program_prefetches` | int64 | 1 |
| `xla_max_cross_program_prefetches` | int | 1 |
| `xla_tpu_llo_compilation_max_retries` | int32 | 15 |
| `xla_tpu_scoped_vmem_limit_kib` | int64 | **-1** (none) |
| `xla_tpu_prefetch_interval_picker_size_override` | int | **-1** (none) |
| `xla_vf_max_vmem_used_by_memory_space_assignment` | int64 | **-1** (none) |
| `xla_gf_max_vmem_used_by_memory_space_assignment` | int64 | **-1** (none) |
| `xla_tpu_max_cmem_used_by_memory_space_assignment` | int64 | **-1** (none) |
| `xla_tpu_sliced_prefetch_max_slices` | int32 | `0xFFFFFFFF` (unset) |
| `xla_tpu_sliced_prefetch_min_bytes` | int64 | -1 (unset) |
| `xla_tpu_sliced_prefetch_preferred_slice_size` | int64 | -1 (unset) |
| `xla_tpu_auto_spmd_partitioning_memory_budget_ratio` | float | 1.1 |

`xla_tpu_msa_inefficient_use_to_copy_ratio`（0.5）是 alt-mem skip threshold：如果 in-alt-mem idle time 与 copy time 的比例超过 0.5，该 buffer **不会**被放入 alternate memory。字节精确验证：default-gen @ `0x1d721c60` 写入 `0x3F000000` = 0.5f。

> **GOTCHA —** alternate-memory byte budget（`Options.max_size_in_bytes`）**不是**从任何 flag 读取的。每个 `*_max_vmem/cmem_used_by_memory_space_assignment` flag 和 `xla_tpu_scoped_vmem_limit_kib` 都默认 `-1`，表示“从 `Target` 推导”。budget 是 VMEM 硬件大小（v7x / 6acc60406 上为 64 MiB = 67,108,864 B，见 [tpu-chip-config.md](../targets/tpu-chip-config.md)）减去任何 scoped reservation，由 [chip-parts-binarypb.md](../targets/chip-parts-binarypb.md) 填充，而不是来自 flag literal。寻找“MSA budget” flag 的重新实现会只找到 `-1`，并且必须回退到 chip config。

### `copies_limit_for_sync_mem_op_conversion` 是计算值，不是 flag

这个边界**不是**顶层 flag。它是 MSA `Options` struct 的派生字段（TextFormat key `extend_async_copies_limit_for_sync_mem_op_conversion:`），在 `ComputeMemorySpaceAssignmentOptions` 内由 outstanding-copy caps 和两个 sync-replacement enables（`xla_msa_enable_sync_copy_replacement`、`xla_msa_enable_sync_slice_replacement`，二者 AutoProto 默认 = unset/AUTO）计算得到。它没有每版本字面值。

> **NOTE —** `copies_limit_for_sync_mem_op_conversion` 的 basic-block 级公式未追踪。它是 outstanding caps（4 或 40）加两个 sync-replacement enables 的函数；由于这些 enables 默认 AUTO/unset，sync→async conversion 实际上被门控关闭，除非被 override。（未追踪 — 明确缺口。）

### Boolean / Tristate global enables

| Flag | 类型 | 默认 |
|---|---|---:|
| `xla_msa_allocate_scoped_memory_at_same_offset` | bool | 1 |
| `xla_tpu_allocate_scoped_vmem_at_same_offset` | bool | 1 |
| `xla_tpu_allocate_scoped_cmem_at_same_offset` | bool | 0 |
| `xla_msa_use_bundle_aware_cost_model` | Tristate | 2 (ENABLED) |
| `xla_msa_cross_program_prefetch_permissive_mode` | bool | 0 |
| `xla_msa_experimental_use_telamalloc` | Tristate | 0 (AUTO) |
| `xla_tpu_vmem_use_telamalloc` / `cmem_use_telamalloc` | bool | 0 / 0 |
| `xla_tpu_msa_use_minimalloc` / `use_tinymalloc` | bool | 0 / 0 |
| `xla_tpu_msa_reduce_scoped_vmem_limit` | bool | 0 |
| `xla_msa_enable_sync_copy_replacement` | AutoProto | 0 (AUTO) |
| `xla_msa_enable_sync_slice_replacement` | AutoProto | 0 (AUTO) |
| `xla_msa_enable_window_prefetch` | AutoProto | 0 (AUTO) |
| `xla_vf_allow_split_vmem` | AutoProto | 0 (AUTO) |
| `xla_msa_expanded_scoped_alternate_memory_mode` | AutoProto | 0 (UNDEFINED) |

> **GOTCHA —** AutoProto 包装的旋钮（`sync_copy`/`sync_slice` replacement、`window_prefetch`、`expanded_scoped_alternate_memory_mode`、`allow_split_vmem`）都默认空/AUTO 实例（0）。“AUTO” 表示“flag 层无意见”：具体 `Target` 可能在未追踪的每版本 C++ 中强制它们 ON。不要把 AUTO=0 读作“功能禁用”；应读作“决策延后”。

---

## 默认值如何存储与物化

### `Flag<T>` 对象布局

每个 MSA 旋钮都是一个 `ABSL_FLAG`，其对象是 `.data` 中 0x60 字节的 `absl::flags_internal::Flag<T>`（FlagImpl）。默认值位于 offset `+0x48`：

```text
absl::flags_internal::Flag<T>   (0x60 bytes, .data)
  +0x00  vtable (FlagImpl, 0x22040f18)
  +0x08  const char* name        → .rodata flag-name string
  +0x10  const char* type_name   → "int64_t" / "float" / "bool" / …
  +0x18  const char* filename    → tpu_compilation_environment.cc
  +0x20  FlagOps<T>              → e.g. FlagOps<long> @ 0xe8cd240
  +0x28  HelpGen function        → AbslFlagHelpGenFor<flag>::NonConst
  +0x30  packed metadata word    (data_guard / value-kind bits)
  +0x38  lazy value sentinel     (-1 / 0xff before first access)
  +0x40  default-kind word       (0)
  +0x48  DEFAULT-VALUE UNION:
           kGenFunc  → AbslFlagDefaultGenFor<flag>::Gen
                       body = `mov dword[rdi], imm32; ret` (float)
                            or `mov byte[rdi], imm8; ret`  (bool/enum)
           kOneWord  → inline int64 literal (the int caps: 4 / 40 / 2 / -1 / 0xFFFFFFFF)

QUIRK — floats 和 bools 使用 +0x48 处的 Gen function（很小的 mov imm; ret body）；integer caps 使用 kOneWord inline literal。这就是为什么 max_outstanding_prefetches 和 max_repacks 在反编译中没有 Gen 符号：它们的默认值是 Flag<T> 对象中的 raw immediate，通过直接读取 +0x48 恢复，而不是通过反汇编 Gen body。扫描 AbslFlagDefaultGenFor* 的重新实现者会找到每个 float ratio，但找不到任何 int caps。

两阶段解析

text

STAGE 1 — materialize defaults (once, lazily)
  GetTpuCompEnvWithDefaultValues()                @ 0x1d73f100
    └─ $_0::operator()()                          @ 0x1d73f1a0   per-field loop
         for each FieldDescriptor in TpuCompilationEnvironment:
           flag = TpuCompEnvReflection::GetFlagForField(field)   @ 0x1d74ad40
           SetFieldFromFlagString(flag, default, env)            @ 0x1d73fcc0
    └─ result: default TpuCompilationEnvironment proto
                (singleton @ 0x22803928, sizeof 0x15e8, guard @ 0x2257ec08)

STAGE 2 — per-version selection (per compile)
  ComputeMemorySpaceAssignmentOptions(Target&, AliasInfo*, HloModule&)  @ 0x12fc1440
    for each MSA Option:                          (55 calls total)
      OverwriteFieldIfNotDefault(                 @ 0x1d73f360
        "xla_<family>_vmem_<opt>",   // family chosen by jump table @ 0xae09ac8
        "xla_msa_<opt>",             // global fallback
        env)
```text

`GetTpuCompEnvWithDefaultValues`（Stage 1）迭代每个 `FieldDescriptor`，通过 reflection 解析匹配的 `CommandLineFlag`，并把 flag 的当前默认值复制到字段中。结果是默认 `TpuCompilationEnvironment` proto：**默认 proto 等于 `absl` flag defaults**，不是硬编码表，也不是 binarypb overlay。process-new-env hook `ProcessNewTpuCompilationEnvironment` @ `0x1d742c80`（通过 `RegisterProcessNewEnvFn` 注册）将这些默认值接入实时 `CompilationEnvironments`。

`ComputeMemorySpaceAssignmentOptions`（Stage 2）是发生每版本选择的地方：对每个 Option，它调用 `OverwriteFieldIfNotDefault(family_flag, msa_fallback, env)`，family 由 `0xae09ac8` 处的 version jump table 选择。反编译显示，对 overlap ratios 和 int caps，所有五个 family prefixes 都与 `xla_msa_*` global 成对出现，总共 55 次 `OverwriteFieldIfNotDefault` 调用。

> **NOTE —** 此构建中**没有 env binarypb overlay**。默认值就是 `absl` flag defaults，到此为止。后续面向 v5e/v6e 的 wheel 可能发布字面的 `TpuCompilationEnvironment` binarypb；此 0.0.40 wheel 没有（filewrapper TOC 只携带 chip_parts/chip_configs/bootloaders/accuracy/route-cache/ICI/CSR-error/tz）。见 [tpu-compilation-environment.md](../config/tpu-compilation-environment.md)。

---

## 工作链路 — 6acc60406 (v5) 上的 `max_overlap`

对 v5（6acc60406）target 上的单个旋钮 `max_overlap_to_mem_size_async_copy_ratio`，端到端如下：

```c
// 1. version → family
tpu_version = *(int32*)(Target + 0x398);              // = 5
family = jump_table_0xae09ac8[5];                     // → gf (shared by Ghostlite + 6acc60406)

// 2. per-version selection
OverwriteFieldIfNotDefault(
    "xla_gf_vmem_max_overlap_to_mem_size_async_copy_ratio",   // family flag
    "xla_msa_max_overlap_to_mem_size_async_copy_ratio",       // global fallback
    env);                                             // inside ComputeMSAOptions @ 0x12fc1440

// 3. the family flag's literal default
AbslFlagDefaultGenForxla_gf_vmem_max_overlap_…::Gen   // @ 0x1d72dfa0
    *(uint32*)rdi = 0x41000000;                       // = float 8.0
// (had this been Jellyfish, jf gen @ 0x1d72cae0 = 0x42000000 = 32.0)

// 4. consumed by the prefetch window
PrefetchIntervalPicker.ctor(...)                      // @ 0x1dcd6b60
    picker[0x80] = min(1.0);
    picker[0x84] = preferred(2.0);
    picker[0x88] = derived(preferred × max=8.0);
    // Begin: latest_prefetch_time = use_time − ⌈min(1.0) × async_copy_elapsed⌉

因此在 v5（6acc60406）target 上，buffer 最多可提前约其自身 copy time 的 8× 被 prefetch；picker 在该边界与使用点之间扫描候选 start times。相同的 8× 适用于 Ghostlite（v4），它共享 gf flags。在 Jellyfish 上，同一旋钮解析为 32×，这是唯一具有非默认 max ratio 的 family。

MSA Config Proto Schema

上游 XLA memory_space_assignment.proto 嵌入在 protodesc_cold 中（FileDescriptorProto @ VA 0xbfe06e0，2,939 B，proto3）。相关消息形状如下（字节精确解码）：

text

message SlicedPrefetchOptions {
  uint32 max_slices = 1;  uint64 min_bytes = 2;
  bool   fail_on_non_alignment_boundary_slice_proposal = 3;
  uint32 all_slice_time_permutations_threshold = 4;
  uint64 preferred_slice_size = 5;
}
message MemoryBoundLoopOptimizerOptions {
  bool  enabled = 1;  float desired_copy_ratio = 2;
  bool  allow_unsatisfied_fully_pipelined_prefetch = 3;
  float min_num_iterations = 4;          // FLOAT
}
enum ExpandedScopedAlternateMemoryMode.Value { UNDEFINED=0; DISABLED=1; ENABLED=2; }
```text

> **GOTCHA —** 由于 `memory_space_assignment.proto`（@ `0xbfe06e0`）和 `tpu_compilation_environment.proto`（@ `0xbfa6060`，137,692 B，1,121 fields）都是 **proto3**，两个 descriptor 都不携带 `default_value` 字符串。数值默认值不在 descriptor 中，而是在 `absl` flag 层（见上）。TCE proto 中的 MSA flag 字段通过 custom field-option extension `#535801365`（`TpuCompEnvFieldOptions{is_used_at_runtime=1, tags=101}`）标记为 `MEMORY_SPACE_ASSIGNMENT (101)`。期望通过解析 descriptor 找到每字段默认值的重新实现者什么也找不到，必须改读 flag objects。

`MemoryBoundLoopOptimizerOptions` 默认值（`desired_copy_ratio`、`min_num_iterations`）是*嵌入消息*，不是 scalar flags；它们的默认值位于 message default-instance 中，此处**未**解码（`xla_tpu_memory_bound_loop_optimizer_options`，TCE field 568；ctor @ `0x1e6c5120`）。（未追踪 — 明确缺口。）

---

## 证据锚点

| Symbol / Anchor | VA | 角色 |
|---|---|---|
| `IsMemorySpaceAssignmentEnabled` | `0x12fc1280` | Version gate + family select |
| `ComputeMemorySpaceAssignmentOptions` | `0x12fc1440` | Per-version selection（55 次 Overwrite 调用） |
| version jump table | `0xae09ac8` | 6 项：v0/1→jf，v2→cmem，v3→vf，v4/5→gf |
| `GetTpuCompEnvWithDefaultValues` | `0x1d73f100` | Default-proto materializer |
| `…::$_0::operator()` | `0x1d73f1a0` | Per-field copy loop |
| default `TpuCompilationEnvironment` | `0x22803928` | singleton（sizeof 0x15e8），guard `0x2257ec08` |
| `TpuCompEnvReflection::GetFlagForField` | `0x1d74ad40` | field → CommandLineFlag |
| `SetFieldFromFlagString` | `0x1d73fcc0` | 将 flag default 复制到字段 |
| `ProcessNewTpuCompilationEnvironment` | `0x1d742c80` | process-new-env hook |
| `OverwriteFieldIfNotDefault` | `0x1d73f360` | resolution primitive |
| `jf max_overlap` Gen (`0x42000000`=32.0) | `0x1d72cae0` | 字节精确 float |
| `gf max_overlap` Gen (`0x41000000`=8.0) | `0x1d72dfa0` | 字节精确 float |
| `vf max_overlap` Gen (`0x41000000`=8.0) | `0x1d72d4a0` | 字节精确 float |
| `jf min_overlap` Gen (`0x3F800000`=1.0) | `0x1d72cc80` | 字节精确 float |
| `jf preferred_overlap` Gen (`0x40000000`=2.0) | `0x1d72cd60` | 字节精确 float |
| `inefficient_use_to_copy_ratio` Gen (`0x3F000000`=0.5) | `0x1d721c60` | 字节精确 float |
| `xla_msa_enable` Gen (byte 2) | `0x1d705500` | 字节精确 Tristate |
| `msa_enable_cross_program_prefetch_freeing` Gen (byte 1) | `0x1d705e00` | 字节精确 bool |
| `msa_use_bundle_aware_cost_model` Gen (byte 2) | `0x1d72f260` | 字节精确 Tristate |
| `zf max_outstanding_prefetches` (=40) | `0x1d72ef40` | forward-stub flag（未门控） |
| `memory_space_assignment.proto` FDP | `0xbfe06e0` | proto3，2,939 B |
| `tpu_compilation_environment.proto` FDP | `0xbfa6060` | proto3，137,692 B，1,121 fields |

---

## 未决项

以下内容未追踪，仍为 LOW/unknown：

- **`copies_limit_for_sync_mem_op_conversion` 公式** — 在 `ComputeMemorySpaceAssignmentOptions` 内由 outstanding caps + sync-replacement enables 计算；basic-block 级推导未恢复。
- **`MemoryBoundLoopOptimizerOptions` 消息默认值** — `desired_copy_ratio` / `min_num_iterations` 位于嵌入消息 default-instance（TCE field 568）中，此处未解码。
- **`zf` 激活路径** — 该 family 的 outstanding caps 存在（=40），但它缺席 v0..v5 jump table；会路由到它的版本（或 `Target` 级 override）未知。
- **Pufferfish (v2) overlap/repack/retry** — v2 路由到 `cmem` 层级（min=1.0/pref=2.0/max=8.0，caps 40，repacks=4，retries=2）；v2 是否还通过 `Target` 级 override 咨询某个 vmem family 未重新推导。
- **AUTO knobs 的每 `Target` overrides** — 任何为特定 `Target` 强制 sync-replacement / window-prefetch / split-vmem knob ON 的每版本 C++ 均未追踪。

---

## 交叉引用

- [msa-overview.md](msa-overview.md) — 第 7 阶段 memory-space-assignment pass；本页提供其每版本数值配置。
- [msa-allocate-segment.md](msa-allocate-segment.md) — 这些默认值参数化的分配主体（outstanding caps、repacks/retries）。
- [msa-reservation-hbm-policy.md](msa-reservation-hbm-policy.md) — 与 chip-config budget 一起设置 `max_size_in_bytes` 的 reservation/HBM policy。
- [compile-phases.md](compile-phases.md) — 第 7 阶段放置；`RunMemorySpaceAssignment` @ `0x12fc3080`。
- [tpu-compilation-environment.md](../config/tpu-compilation-environment.md) — 这些 flags 填充的 1,121 字段 TCE proto（无 binarypb overlay）。
- [tce-field-offsets-defaults.md](../config/tce-field-offsets-defaults.md) — field#→offset 映射，以及所有 TCE 字段共享的 ABSL-flag default 机制。
- [tpu-chip-config.md](../targets/tpu-chip-config.md) — 提供 VMEM byte budget（`max_size_in_bytes`）的每版本 chip config，因为 MSA byte-cap flags 默认 `-1`。
- [chip-parts-binarypb.md](../targets/chip-parts-binarypb.md) — budget 推导自的 `chip_parts` 资源。
- [tpu-version-codename-matrix.md](../targets/tpu-version-codename-matrix.md) — jf/cmem/vf/gf family selection 背后的 `tpu_version` ↔ codename 映射。

MSA 每版本默认值 ​

摘要 ​

代号 Family 与版本门控 ​

目的 ​

Family 前缀 ​

算法 ​

两阶段解析 ​

MSA Config Proto Schema ​

MSA 每版本默认值

摘要

代号 Family 与版本门控

目的

Family 前缀

算法

两阶段解析

MSA Config Proto Schema