常量映射器

地址适用于 libtpu-0.0.40-cp314 轮中的 libtpu.so。其他版本有所不同。

摘要

ConstantMapper 是水母集体降低的编译时常量池：降低到 ICI 程序的每个集体 HLO(all-reduce、all-gather、all-to-all、collective-permute)都需要一些静态表 - 副本/序号表、ND 环邻居表、每次传输路由表、二项式蝶形表、所有到所有屏障成员资格表，其中每一个都注册到一个由小整数 Type 枚举键控的 ConstantMapper 中，并随后从中获取。该常量池的相反编号是 运行时集合 id：TPU 核心在编译时不知道自己的副本/分区等级，因此每个等级在运行时作为单个 U32 标量从固定 SMEM 字读取，该固定 SMEM 字的偏移量本身在芯片配置中的 Target::Init 处保留。此页面拥有 Type 枚举、GenerateConstants → GetConstantTables 实现路径以及 SMEM-id 读取。

该机制不是全局枚举索引表。 ConstantMapper 是一个每降低指令对象，保存由 Type int 键控的 Abseil SwissTable(FlatHashMapPolicy<ConstantMapper::Type, StatusOr<shared_ptr<const vector<int>>>>，从损坏的实例化中确认)。相同的 Type 值在不同的发射器中携带不同的表内容 - Type 3 是 AllReduce 中的平面复制表，但在 AllGather 中是 ND 环表 - 并且这正是无冲突的，因为每个降低的集体拥有自己的映射器，由每个 HLO 工厂 (GetConstantFnForCollective) 分发。第二个命名空间 MeshNDInfo::MeshDim 共享相同的整数：MeshDim AddConstant/GetConstant 重载将值转发到 Type 重载不变，因此网格轴 0/1/2 是 Type 0/1/2。

对于重新实现者来说，合约由三个部分组成：(1) 12 值 Type 枚举 (0..0xb) 以及哪个表构建器将每个标签提供给哪个发射器； (2) 注册/回读路径 — GenerateConstants 填充映射器，降低主体由 GetConstant(Type) 获取，AllToAll 屏障专门将类型 8/9/0xa 通过 GetConstantTables 作为 (InfoTable, InfoTable, optional<InfoTable>) 三元组 — 包括静态-Literal-与动态-vector<int> 运营商选择； (3) 运行时 ID 读取：GetReplicaId/GetPartitionId 从 Target+0x6f8/+0x700 发出 U32 Sld(标量负载)，从芯片配置的用户保留 SMEM 区域保留字偏移。全对全表(类型8/9/0xa)的表内容归AllToAll 表所有； AllGather ND 环的 ND 环表(类型 0/1/2、3)。该页面拥有枚举、载体机制和读取的 id。

重新实现，合约为：

Type 枚举 — 12 个标签 (0..0xb)、MeshDim ≡ Type 别名、类型 3/4 的每发射器过载以及一个跨集合标签(类型 5 = 路由调度，由AllToAll/CollectivePermute/AllGather)。
物化路径 — GetConstantFnForCollective(每个HLO工厂)→ GenerateConstants(AddConstant(Type, …))→ GetConstant(Type) / GetConstantTables； SwissTable 存储以及静态与动态载体门。
运行时集合 ID 读取 — GetReplicaId/GetPartitionId = U32 Sld 位于目标驻留字偏移处； Target::Init 保留芯片配置用户保留的 SMEM 区域的偏移量； partition_count==1 折叠至 0。


映射器类型	`xla::jellyfish::ConstantMapper` — 每个降低的指令对象
存储	Abseil SwissTable `FlatHashMap<ConstantMapper::Type, StatusOr<shared_ptr<const vector<int>>>>`
密钥空间	`ConstantMapper::Type` int 0..0xb(12 个标签)； `MeshNDInfo::MeshDim` 别名相同的整数
添加(文字)	`ConstantMapper::AddConstant(Type, StatusOr<Literal>)` @ `0x1c885ce0`
添加(矢量)	`ConstantMapper::AddConstant(Type, StatusOr<vector<int>>)` @ `0x1c886300`
获取	`ConstantMapper::GetConstant(Type)` @ `0x1c886b00`； `HasConstant(Type)` @ `0x1c886920`
工厂	`GetConstantFnForCollective` @ `0x10c46f60`(每个 HLO 闭合 `$_0..$_3`)
屏障读回	`GetConstantTables` @ `0x10f07860` → `(InfoTable, InfoTable, optional<InfoTable>)`
副本 ID 读取	`net_util::GetReplicaId` @ `0x1c69a440` → U32 `Sld` 在 `Target+0x6f8`
分区 ID 读取	`net_util::GetPartitionId` @ `0x1c69a4a0` → U32 `Sld` 在 `Target+0x700`
偏移保留	`Target::Init` @ `0x1d60fc20` 通过 `GetUserReservedSmemBlock` @ `0x1d613b20`

类型枚举

用途

ConstantMapper::Type 是用作 SwissTable 键的平面整数枚举，而不是全局表索引。它标记了集体发射器所需的每个静态常数，以便降低主体可以在发射期间通过标签取回每个静态常数。有 12 个正在使用的标签，从 0 到 0xb，通过在每个 AddConstant/GetConstant/HasConstant 调用站点读取 Type 立即数并将其与生成值的表构建器配对来建立。

机制 — 每个实例 SwissTable，`MeshDim` ≡ `Type`

每个降低的集体 HLO 都拥有一个 ConstantMapper。 GenerateConstants(hlo, target, topo, region) 构建它所需的常量 - 副本/分区/路由/二项表，作为 xla::Literal R1 int 常量或 vector<int> - 并通过 AddConstant 将每个常量注册到 Type 下。随后，下降体由 GetConstant(Type) 取回。两个AddConstant(Type, …)重载(Literal @ 0x1c885ce0、vector<int> @ 0x1c886300)和GetConstant(Type) @ 0x1c886b00是整个API表面； HasConstant(Type) @ 0x1c886920 是一个可选标签的存在探针。

MeshNDInfo::MeshDim 共享整数命名空间。 MeshDim 重载没有自己的存储 - 它们将网格轴索引不变地转发到 Type 重载：

// ConstantMapper::AddConstant(MeshDim, StatusOr<Literal>)  — 0x1c886260
function AddConstant_MeshDim(this, mesh_dim /*esi*/, statusor_literal):
    // unwrap the StatusOr<Literal> into a bare Literal, then forward
    // mesh_dim (esi) UNCHANGED into the Type overload:
    return AddConstant_Type(this, /*Type=*/mesh_dim, literal)   // tail-call 0x1c885ce0

// ConstantMapper::GetConstant(MeshDim)  — 0x1c887560   (// attributes: thunk)
function GetConstant_MeshDim(this, mesh_dim):
    return GetConstant_Type(this, mesh_dim)                     // jmp 0x1c886b00
```text

`GetConstant(MeshDim)` 是一个跳转到 `GetConstant(Type)` 的单指令 thunk； `AddConstant(MeshDim, vector<int>)` @ `0x1c886620` 以同样的方式转发到 `AddConstant(Type, vector<int>)`。因此，每个网格轴 ND 环表将注册并读回为类型 0/1/2 — 每个轴的 `MeshDim`。

> **QUIRK —** `Type` 和 `MeshDim` 是相同的整数，但不同的 C++ 类型，因此为它们提供单独的哈希映射的重新实现将在稍后查找 `Type 1` 时默默地找不到 `MeshDim 1` 下添加的常量。它们必须为相同的密钥空间设置别名，就像二进制文件将 `MeshDim` 转发到 `Type` 重载而无需重新映射一样。

### 12个标签(0..0xb)

每行都固定提供 `AddConstant(Type=k)` 调用的表构建器(因此表 *内容* 被锚定)和读回它的使用者。重载/不透明标签的英文角色名称来自生产构建器，而不是从枚举器描述符中读取 - 标记为 `HIGH`，其中生产者是唯一的证据。

| 型 | 载体 | 表构建器(生产者) | 内容/角色 |
|---|---|---|---|
| 0 | `MeshDim 0` | `CreateStaticNDRingReplicaInfoTable` / `CreateNDRingReplicaInfoTable` | ND 环复制工作台，网格 **轴 0** |
| 1 | `MeshDim 1` | (相同，轴1) | ND 环复制工作台，网格 **轴 1** |
| 2 | `MeshDim 2` | (相同，轴2) | ND 环复制工作台，网格 **轴 2** |
| 3 | `Type`(重载) | AllReduce：`CreateReplicaInfoTable[ForLimitedIciRouting]`；全部聚集：`CreateStaticNDRingReplicaInfoTable` | 平面组内副本/序数表 (AR) **或** ND 环表 (AG) |
| 4 | `Type`(重载) | `CreateReplicaInfoTableForLimitedIciRouting` / `CreateStaticReplicaInfoTableForLimitedIciRouting` | Limited-ICI-路由副本表/路由表索引 |
| 5 | `Type`(**跨集体**) | `net_router::CreateRoutingScheduleLiteral` / `CreateAllToAllRoutingScheduleTable` | 每次传输 ICI **路线时间表**文字 |
| 6 | `Type` | `net_util::CreateNDRingReplicaInfoTable`(AllReduce) | 环AllReduce ND重新排序表 |
| 7 | `Type` (+`Status`) | `CreateBinomialReplicaInfoTable` | 二项式递归加倍蝶形表 |
| 8 | `Type` | `GenerateAllToAllTables`(表A) | AllToAll 屏障成员资格表 A |
| 9 | `Type` | `GenerateAllToAllTables`(表B) | AllToAll 屏障成员资格表 B |
| 0xa | `Type`(`HasConstant`-门控) | `GenerateAllToAllTables`(表C) | AllToAll **可选** 成员资格表 C |
| 0xb | `Type` | `CreateCollectivePermuteTransfers` | CollectivePermute传输/源-目标表 |

三个结构事实驱动该表：

- **类型 5 是唯一的跨集合标签。** 三个路由驱动的发射器 — AllToAll、CollectivePermute 和 AllGather(显式路由路径) — 生成*并*使用它；它携带 `net_router` 的路由调度文字——每次传输 ICI 路由程序。 (AllReduce **不**触及类型 5：它的环/二项式算法使用类型 3/4/6/7 代替 — 字节确认，`AllReduceEmitter::GenerateConstants` 中没有 `AddConstant(…, 5, …)`。)其内部 int 布局此处未解码(请参阅 [创建路由计划](../routing/create-routing-schedule.md) / [路由表生成](../routing/route-table-generation.md))。
- **类型 3 和 4 是重载槽索引。** 同一标签在 AllReduce 中携带平面副本表，在 AllGather 中携带 ND 环/静态表。不存在冲突，因为映射器是按降低的指令进行的(请参阅 [§ 物化路径](#the-materialization-path))。
- **类型 7 同时采用表和 `Status`。** 在非二项式可行路径上，AllReduce 发射器执行 `AddConstant(Type=7, Status)`；后来的 `GetConstant(7)` 然后显示该错误。这是二项式表标签 — 请参阅 [二项式递归加倍](binomial-recursive-doubling.md)。

> **注意 —** 类型 8/9/0xa *是 AllToAll 副本组内屏障的 `(InfoTable A, InfoTable B, optional<InfoTable> C)`，由 `GetConstantTables` @ `0x10f07860` 作为三元组读回。表内容、两条填充路径、索引算法归[AllToAll 表](alltoall-tables.md)所有；该页面仅拥有标签分配和读回路径。

### 生产者和消费者

通过在每个调用站点立即读取 `Type`/`MeshDim` 来重建枚举。生产者/消费者分裂(每个 `AddConstant` 都是生产者，每个 `GetConstant`/`HasConstant` 都是消费者)：

| 发射器 | 制作人 — `GenerateConstants` | 添加标签 | 消费者 |
|---|---|---|---|
| AllToAll | `AllToAllEmitterBase::GenerateConstants` @ `0x10f089a0` | 8, 9, 0xa, 5 | `CalculateWithLimitedIciRouting` (5)、`GetConstantTables` (8/9/0xa) |
| 集体置换 | `CollectivePermuteEmitter::GenerateConstants` @ `0x1346ff60` | 0xb, 5 | `EmitForLimitedIciRouting` (5)、`Emit` (0xb) |
| AllReduce | `AllReduceEmitter::GenerateConstants` @ `0x1373cb60` | 3, 4, 6, 7(+状态); `MeshDim 0/1`(单独的跨模块 ARS 映射器) | `GetRingLocation` (3)、`GetRingLocationWithReordering` (4)、`EmitAllReduceFusion`/`ConstructAsyncFusionEmitter` (6)、`BuildStrategyForCrossModuleARS` (0,1)、`BinomialGroupData $_1` (7) |
| AllGather | `AllGatherEmitter::GenerateConstants` @ `0x13801be0` | `MeshDim 0/1/2`, 3, 4, 5 | `InitDim` (3/4/`MeshDim`)、`EmitAllGatherWithExplicitRouting` (5) |

> **明白了 —** 避免“每个集合都存在所有 12 个标签”的重新实现是错误的。每个发射器仅填充其算法所需的标签； AllToAll 永远不会添加类型 3，AllReduce 永远不会添加类型 8。`Type` 空间是发射器上的*联合*，而不是每个实例模式。仅添加所选算法产生的内容，仅获取其消耗的内容。

---

## 物化路径 {#the-materialization-path}

### 用途

这是将 HLO 集合转换为填充的 `ConstantMapper` 的调度和生命周期，然后在发射期间读回表：一个工厂通过操作码选择正确的 `GenerateConstants`，该函数填充新的映射器，并且降低主体(或屏障读回)通过标签获取。

### 入口点

```text
GetConstantFnForCollective (0x10c46f60)        ── HLO-opcode → factory closure
  └─ $_0 / $_1 / $_2 / $_3                      ── forward into one GenerateConstants
       └─ AllReduce/AllGather/AllToAll/CollectivePermute::GenerateConstants
            └─ AddConstant(Type, …)              ── populate the per-instance SwissTable
                 ⋮
       (lowering body)  GetConstant(Type)        ── 0x1c886b00, fetch by tag
       (AllToAll barrier) GetConstantTables      ── 0x10f07860, read 8/9/0xa as a triple

算法 — 工厂、填充、获取

GetConstantFnForCollective 返回一个 std::function<StatusOr<unique_ptr<ConstantMapper>>(HloInstruction*)> ，其目标是四个 $_0..$_3 闭包之一(在反编译体中确认：该函数安装 __policy_func::__call_func<…::$_0> 等)。每个闭合件转发到匹配发射器的GenerateConstants。因此，每个集体指令有一个映射器，由操作码的 GenerateConstants 延迟构建：

function GetConstantFnForCollective(hlo, target, topo, region):   // 0x10c46f60
    switch opcode_class_of(hlo):
        all-reduce:        bind closure $_? → AllReduceEmitter::GenerateConstants
        all-gather:        bind closure $_? → AllGatherEmitter::GenerateConstants
        all-to-all:        bind closure $_? → AllToAllEmitterBase::GenerateConstants
        collective-permute:bind closure $_? → CollectivePermuteEmitter::GenerateConstants
    return std::function<StatusOr<unique_ptr<ConstantMapper>>(HloInstruction*)>

function GenerateConstants(hlo, target, topo, region):   // per-emitter
    mapper = new ConstantMapper
    for each table the algorithm needs:
        value = <table builder>(...)              // CreateReplicaInfoTable / GenerateAllToAllTables / …
        AddConstant(mapper, /*Type=*/tag, value)  // 0x1c885ce0 (Literal) or 0x1c886300 (vector)
    return mapper
```text

`AddConstant(Type, Literal)` @ `0x1c885ce0` 读取 `Type` int (`esi`→`r15d`) 并存储由 int 键控的 448 字节 (`0x1c0`) 记录`record[+0]` — 在支持 SwissTable 的槽向量上进行线性扫描/插入(`cmp record[+0]==Type`，跨步 `imul $0x1c0`)。 `GetConstant(Type)` @ `0x1c886b00` 是哈希查找：它读取 `Type` (`edx`→`r15d`)，对其进行 CRC32 哈希，对控制字节进行分组探测，并将候选密钥与请求的 `Type` 进行比较。

### 静态与动态载流子门

`AddConstant` 有两个重载，生产者为每个表选择是否烘焙**静态** R1 `Literal`(`AddConstant(Type, StatusOr<Literal>)` @ `0x1c885ce0`)或携带**动态** `vector<int>`(`AddConstant(Type, StatusOr<vector<int>>)` @ `0x1c886300`)。静态 `Literal` 是降低程序直接从物化 R1 缓冲区读取的常量；动态 `vector<int>` 在运行时具体化。同一标签的生产者调用站点通常出现两次 - 一次带有 `vec` 参数，一次带有 `lit` 参数 - 这是载体分支：当表在编译时完全已知时，发射器选择静态 `Literal` 形式，否则选择动态 `vector` 形式。所有对所有运营商的选择在 [AllToAll 表](alltoall-tables.md) 上进行了详细解码。

> **注意 —** SwissTable 值类型是 `StatusOr<shared_ptr<const vector<int>>>`(来自损坏的 `FlatHashMapPolicy` 实例化)。 `Literal` 过载仍然通过相同的地图进行路由；静态与动态的区别在于*降低对获取的常量的作用*(读取烘焙的 R1 缓冲区与物化)，而不是在两个单独的映射中。

### 屏障读回 — `GetConstantTables`

AllToAll 副本组内屏障不直接调用 `GetConstant`；它通过 `GetConstantTables`，读取三个成员资格标签并将它们作为屏障的 `InfoTable` 三元组返回：

```c
function GetConstantTables(hlo, mapper):           // 0x10f07860
    A = GetConstant(mapper, /*Type=*/8)            // 0x1c886b00 — barrier InfoTable A
    if A is error: return A
    B = GetConstant(mapper, /*Type=*/9)            //            — barrier InfoTable B
    if B is error: return B                        // AddSourceLocation: all_to_all_emitter_base.cc:323
    C = optional{}
    if HasConstant(mapper, /*Type=*/0xa):          // 0x1c886920 — HasConstant(a2, 10)
        C = GetConstant(mapper, /*Type=*/0xa)      //            — optional InfoTable C
    return tuple<InfoTable, InfoTable, optional<InfoTable>>(A, B, C)

反编译主体显示了三个 GetConstant 调用和第三个调用的 HasConstant(a2, 10) 门(10 = 0xa)，确认了读取侧门表 C 的存在与 GenerateAllToAllTables 门的构造完全相同。返回的三元组直接进入屏障启动参数寄存器。

功能图

功能	地址	角色
`ConstantMapper::AddConstant(Type, StatusOr<Literal>)`	`0x1c885ce0`	静态-`Literal`注册商；由 `record[+0]==Type` 键入的 448 字节记录
`ConstantMapper::AddConstant(Type, StatusOr<vector<int>>)`	`0x1c886300`	动态-`vector<int>`注册商
`ConstantMapper::AddConstant(Type, Status)`	`0x1c8866c0`	错误载体注册器(类型 7 不可行路径)
`ConstantMapper::GetConstant(Type)`	`0x1c886b00`	SwissTable 查找(`Type` 的 CRC32 哈希)
`ConstantMapper::HasConstant(Type)`	`0x1c886920`	存在探头(门类型 0xa)
`ConstantMapper::AddConstant(MeshDim, …)`	`0x1c886260` / `0x1c886620`	`MeshDim`→`Type` 转发器(esi 不变)
`ConstantMapper::GetConstant(MeshDim)`	`0x1c887560`	重击 `jmp 0x1c886b00`
`GetConstantFnForCollective`	`0x10c46f60`	每HLO工厂；关闭 `$_0..$_3`
`GetConstantTables`	`0x10f07860`	读取类型 8/9/0xa → `InfoTable` 三重

QUIRK — Type-3/Type-4 重载仅是无冲突的，因为映射器是按降低的指令进行的。 GetConstantFnForCollective 每个 HLO 返回一个新的闭包是字节确认的；这为每条指令生成一个映射器是结构基础，并且是 HIGH，而不是 CERTAIN — 跨指令共享一个映射器的重新实现者会将 AllReduce 的平面类型 3 别名为 AllGather 的 ND 环类型 3。

运行时集体 ID 读取

用途

集体发射器需要知道“我是哪个等级”——它的 replica_id 和 partition_id——来索引上面的静态表(二项式表索引是 replica_id*8 + col；(replica_id, partition_id) 上的平面屏障读取器键)。该 id 不是编译时常量：每个核心在运行时从固定的 SMEM 字中读取自己的 id 作为单个 U32 标量。编译器只发出负载；主机运行时/TPU 固件在程序运行之前将每个核心的 ID 存入保留字中。

算法 — 来自目标驻留字偏移的一个 U32 `Sld`

function net_util::GetReplicaId(b /*LloRegionBuilder*/):       // 0x1c69a440
    word_off = b.target()->ReplicaIdLocationWordOffset()       // Target+0x6f8
    ptr      = b.SmemWordImmPtr(word_off, "replica id location")// ImmPtr(off, U32, MS=kSmem)
    return     b.Sld(ptr, /*pred=*/nullptr)                    // CreateScalarLoad → LloValue

function net_util::GetPartitionId(b):                          // 0x1c69a4a0  (identical shape)
    word_off = b.target()->PartitionIdLocationWordOffset()     // Target+0x700
    ptr      = b.SmemWordImmPtr(word_off, "partition id location")
    return     b.Sld(ptr, nullptr)
```text

两者都反编译为上面的三调用体。注释字符串 `"replica id location"` (len 0x13) 和 `"partition id location"` (len 0x15) 是文字参数。

`SmemWordImmPtr` @ `0x1d516880` 构建指针：它断言 `target().SmemWordSizeBytes() == sizeof(uint32_t)` (带有文件字符串 `llo_region_builder.cc` 的 `LloCheckForFailure` 保护)，然后 `MakeValidatedShape(8, …)` — PrimitiveType **8 = U32**，变暗`[]` — 和 `ImmPtr(offset, shape, MemorySpace=5=kSmem, annotation)`。所以 id 是 SMEM 中的标量 U32。

`Sld` @ `0x1d516a20` 在发出负载之前验证操作数的内存空间：

```c
function Sld(b, ptr, pred):                                    // 0x1d516a20
    ms = (ptr->flags_byte[0xb] >> 2) & 0x1F                    // memory-space field
    if (ms - 9) >= 2 && ms != 5:                               // accept kSmem(5) or sflag tiers 9..10
        fail("...llo_region_builder.cc:5365")
    inst = LloInstruction::CreateScalarLoad(ptr, pred, region) // 0x1d516a54
    return b.AppendInstruction(inst)                           // 0x1d516a61

接受的集是 MemorySpace 5 (kSmem) 或 SFLAG 级层 9..10 — 有关内存空间分类，请参阅 SMEM 标量存储器。读取的id始终使用kSmem。

HLO 操作和 `partition_count==1` 折叠

HLO replica-id / partition-id 操作直接降低到这些读数：

LoweringEmitter::HandleReplicaId @ 0x10c34260 → GetReplicaId (@ 0x10c34306)。
LoweringEmitter::HandlePartitionId @ 0x10c33940 → GetPartitionId (@ 0x10c339e6)。

集体发射器通过折叠单分区情况的助手调用读取：

function collective_lowering_utils::GetPartitionId(b, hlo):    // 0x13819500
    if hlo->GetModule()->config().partition_count() == 1:      // module-config[+0x178] == 1
        return b.SimmS32(0)                                     // compile-time constant 0
    else:
        return net_util::GetPartitionId(b)                     // the SMEM read (tail-jump)
```text

当没有模型并行性 (`partition_count == 1`) 时，分区 id 是编译时 `0` 并且不会发出 SMEM 负载；否则它尾部跳转到 SMEM 读取。 (已确认：`module->config[+376] == 1`→`SimmS32(0)`，否则为`net_util::GetPartitionId`。)

> **明白了 —** 这种折叠没有 `replica-id` 类似物。 `HandleReplicaId`始终发出SMEM负载；在单分区情况下，只有 `partition_id` 被折叠为 0。将折叠镜像到replica-id 上的重新实现会不断折叠错误的id。

### 谁读取运行时id

`GetReplicaId` @ `0x1c69a440` 有12个调用点； `GetPartitionId` @ `0x1c69a4a0` 有 8 个。每个集体发射器从这些中读取 `(replica_id, partition_id)`。著名的消费者：`HandleReplicaId`/`HandlePartitionId`(HLO 操作)、`AllToAllEmitter::Init` / `RaggedAllToAllEmitter::Init`、`CollectivePermuteEmitter` 构造函数和屏障、`AllReduceEmitter`/`AllGatherEmitter` 构造函数和 `LoadBinomialReplicaInfoTable` @ `0x1375fca0` — 其表索引正是这个 `replica_id`(参见 [二项式递归加倍](binomial-recursive-doubling.md))。

### 功能图

| 功能 | 地址 | 角色 |
|---|---|---|
| `net_util::GetReplicaId` | `0x1c69a440` | U32 `Sld` 位于 `Target+0x6f8`，注释“副本 id 位置” |
| `net_util::GetPartitionId` | `0x1c69a4a0` | U32 `Sld` 位于 `Target+0x700`，注释“分区 id 位置” |
| `LloRegionBuilder::SmemWordImmPtr` | `0x1d516880` | 构建`ImmPtr(off, U32, MS=kSmem)`；断言字长 == 4 |
| `LloRegionBuilder::Sld` | `0x1d516a20` | MS 验证然后 `CreateScalarLoad` + `AppendInstruction` |
| `LoweringEmitter::HandleReplicaId` | `0x10c34260` | HLO `replica-id` 操作 → `GetReplicaId` |
| `LoweringEmitter::HandlePartitionId` | `0x10c33940` | HLO `partition-id` 操作 → `GetPartitionId` |
| `collective_lowering_utils::GetPartitionId` | `0x13819500` | `partition_count==1` → `SimmS32(0)`，否则 `net_util::GetPartitionId` |

---

## SMEM字偏移保留

### 用途

id 读取使用的字偏移量(`Target+0x6f8`、`+0x700`，...)不是硬编码常量。它们是从 `Target` 对象读取的普通字段，并且这些字段从芯片配置的用户保留的 SMEM 区域填充到 `Target::Init` 中。相对布局(区域中哪个单词持有哪个id)是固定的；绝对偏移量是相对于区域基数的，并根据代号/拓扑而变化。

### 目标 SMEM-id 字段集群 (`Target+0x6c0 .. +0x710`)

每个访问器都是单行字段读取(`mov <off>(%rdi),%rax; ret`)；这些字段是 `long`(8 字节)SMEM 字索引。

| 目标关闭 | 访问器 (@VMA) | 基准 |
|---|---|---|
| `0x6c0` | `OutfeedBasePtrWordOffset` @ `0x1d617bc0` | 出料底座指针 |
| `0x6c8` | `OutfeedProducerHostSyncFlagNumberWordOffset` @ `0x1d617be0` | 输出生产者主机 sflag # |
| `0x6d0` | `CachedOutfeedProducerOffsetWordOffset` @ `0x1d617c00` | 缓存出料生产者偏移量 |
| `0x6d8` | `InfeedPtrLocationWordOffset` @ `0x1d617ba0` | 进给指针 |
| `0x6e0` | `ChipIdLocationWordOffset` @ `0x1d617c20` | 芯片 ID |
| `0x6e8` | `CoreIndexLocationWordOffset` @ `0x1d617c40` | 核心指数 |
| `0x6f0` | `PhysicalChipBoundsLocationWordOffset` @ `0x1d617c60` | 物理芯片边界 |
| `0x6f8` | `ReplicaIdLocationWordOffset` @ `0x1d617c80` | **副本 ID** |
| `0x700` | `PartitionIdLocationWordOffset` @ `0x1d617ca0` | **分区 ID** |
| `0x708` | `SliceIdLocationWordOffset` @ `0x1d617cc0` | 切片 ID |
| `0x710` | `SubsliceOriginLocationWordOffset` @ `0x1d617ce0` | 子切片起源 |

### 算法 — `Target::Init` 从芯片配置中保留

```c
function Target::Init(topology, …, Target* out, …):            // 0x1d60fc20
    reservation = chip_config.GetMemoryReservation()           // 0x20afcf00
    // identity ids come from the TOP user-reserved SMEM region:
    out[+0x6f8] = GetUserReservedSmemBlock(reservation, /*type=*/1, /*top=*/1).word_offset // replica_id
    out[+0x700] = GetUserReservedSmemBlock(reservation, /*type=*/2, /*top=*/1).word_offset // partition_id
    out[+0x708] = GetUserReservedSmemBlock(reservation, /*type=*/3, /*top=*/1).word_offset // slice_id
    out[+0x710] = GetUserReservedSmemBlock(reservation, /*type=*/7, …).word_offset         // subslice_origin
    // +0x6c0..+0x6f0 from TpuMemoryReservation::GetRegionForType / TpuChipConfig::GetUserStack

function Target::GetUserReservedSmemBlock(res, type, top):      // 0x1d613b20
    region = TpuMemoryReservation::GetUserRegion(res)
    blocks = top ? kBlocksTop /*0xb53c180*/ : kBlocksBottom /*0xb53c230*/
    entry  = blocks[ index matching entry.type[+0] == type ]   // 24-byte entries
    if top:    abs_word = region.base + res.base - (entry.offset[+8] + entry.num_words[+0x10])  // counts down
    else:      abs_word = entry.offset[+8] + res.base                                            // counts up
    return { word_offset=abs_word, num_words=entry.num_words, type=entry.type }

GetUserReservedSmemBlock 将请求的 type 与静态 24 字节条目表 (entry = {int type[+0], int pad[+4], long word_offset[+8], long num_words[+0x10]}) 进行匹配，通过 top 布尔值选择 TOP 或 BOTTOM kBlocks 表，然后计算绝对 SMEM 字偏移量。

两个保留块布局表：

text

GetUserReservedTopSmemBlocks::kBlocks @ 0xb53c180   (7 entries, each 1 word)
  type 0 @ w0 | type 1 (replica_id) @ w1 | type 2 (partition_id) @ w2 | type 3 (slice_id) @ w3
  type 4 @ w4 | type 6 @ w5 | type 7 (subslice_origin) @ w6

GetUserReservedBottomSmemBlocks::kBlocks @ 0xb53c230 (3 entries)
  type 5 @ w0 (0x25=37 words) | type 8 @ w0x25 (1 word) | type 9 @ w0x26 (1 word)
```text

因此，`replica_id`/`partition_id`/`slice_id` 在 TOP 区域的相对字 1/2/3 处占用三个连续的 1 字槽。它们的绝对偏移量是相对于区域基数的，并且取决于每个核心的 `TpuMemoryReservation`(因此它根据代号/拓扑而变化，但相对布局是固定的)。

> **注意 —** `UserReservedSmemType` 枚举器名称(1=replica_id、2=partition_id、3=slice_id、7=subslice_origin)源自 `Target::Init` 存储目标(每个 `GetUserReservedSmemBlock(type=k)` 结果存储到匹配的`…LocationWordOffset` 字段)，不是从枚举描述符中读取的 — `HIGH` 置信度。 `type` 整数和 `kBlocks` 字偏移量是字节确认的。

### 编译/运行时拆分

> **GOTCHA —** libtpu 仅*保留插槽* (`Target::Init`) 并*发出读取* (`GetReplicaId`)。它从不将值写入保留的 SMEM 字。主机运行时/TPU固件在内核启动之前将每个内核的`(replica_id, partition_id, slice_id)`存入这些字中；该写入器位于 `libtpu.so` 之外，是从只读使用中推断出来的，此处未观察到。期望编译器发出 id 存储的重新实现将找不到这样的存储。

### 数据路径，端到端

| 舞台 | 功能(VMA) | 输出 |
|---|---|---|
| 保留id SMEM字(编译) | `Target::Init` @ `0x1d60fc20` | `Target+0x6f8`/`+0x700`/`+0x708` 字偏移 |
| ↳ 插槽查找 | `GetUserReservedSmemBlock` @ `0x1d613b20` (`kBlocks`) | 类型 → 区域相对字偏移 |
| 发出 id 读取(降低) | `GetReplicaId` @ `0x1c69a440` / `GetPartitionId` @ `0x1c69a4a0` | `LloValue`(U32 ID) |
| ↳ 构建 SMEM ptr | `SmemWordImmPtr` @ `0x1d516880` | `ImmPtr(U32, MS=kSmem)` 为保留字 |
| ↳ 标量负载 | `Sld` @ `0x1d516a20` (`CreateScalarLoad`) | 每核 `replica_id`/`partition_id` |
| HLO 副本 ID / 分区 ID 操作 | `HandleReplicaId` @ `0x10c34260` / `HandlePartitionId` @ `0x10c33940` | id `LloValue` |

---

## 这里没有解码什么

- **5型内部布局。** 确认为跨集体航线班次承运人；它的 int 编码(每传输路由程序)由路由子系统([创建路由计划](../routing/create-routing-schedule.md)、[路由表生成](../routing/route-table-generation.md))拥有。
- **类型 8/9/0xa 和 0xb 表内容。** 通过 `AddConstant` 调用站点固定到其构建者；每个元素的语义由 [AllToAll 表](alltoall-tables.md) (8/9/0xa) 和 CollectivePermute 发射器 (0xb) 所有。
- **`ConstantMapper::Type` (`Type::k…`) 和 `UserReservedSmemType` 的符号枚举器名称**：标签*整数*、它们的*生产者*和类型→字偏移 `kBlocks` 表是字节确认的；未提取枚举自己的描述符/符号表，因此英文角色名称来自生产者和商店目标。
- **每指令一个映射器**(类型 3/4 重载无冲突的基础)：由返回每 HLO 工厂的 `GetConstantFnForCollective` 支持，未单独证明每个降低的指令恰好是一个映射器。
- **运行时/固件 id 写入器**，将每个核心 id 存入保留字中 — 在 `libtpu.so` 之外。

---

## 交叉引用

- [AllToAll 表](alltoall-tables.md) — 拥有类型 8/9/0xa(障碍成员资格表)的内容以及全对全情况的静态与动态载波解码
- [AllGather ND 环](allgather-nd-ring.md) — 拥有 ND 环复制表(类型 0/1/2、3)以及通过 `GetConstant(MeshDim)` 消耗它们的 `InitDim` 轴行走
- [二项式递归加倍](binomial-recursive-doubling.md) — 7 型二项式蝶形表；它的索引是 `GetReplicaId` SMEM 阅读记录在这里
- [SelectND策略](strategy-nd-picker.md) — 决定运行哪个 `GenerateConstants` 以及填充哪些标签的策略决策
- [On-Pod 集体 — 剖面图](overview.md) — 此常量池服务的集体降低管道
- [创建路由计划](../routing/create-routing-schedule.md) — 跨集体类型 5 路线调度文字生成器
- [路由表生成](../routing/route-table-generation.md) — Type 5 背后的路由表生成
- [SMEM 标量存储器](../memory/smem-scalar-memory.md) — `Sld` 验证的内存空间分类 (`kSmem`=5) 和 id 读取使用的标量加载模型

常量映射器 ​

摘要 ​

类型枚举 ​

用途 ​

机制 — 每个实例 SwissTable，MeshDim ≡ Type ​

算法 — 工厂、填充、获取 ​

功能图 ​

运行时集体 ID 读取 ​