Route-Table 生成

二进制： extracted/libtpu-0.0.40-cp314-cp314-manylinux_2_31_x86_64/libtpu/libtpu.so（build-id 89edbbe81c5b328a958fe628a9f2207d，构建 libtpu_lts_20260413_b_RC00；.text VMA == 文件偏移 0xe63c000，.rodata VMA == 文件偏移 0x84a0000）。 状态： 可重新实现级 · 证据等级： Confirmed (byte-anchored) — DmaDestinationRoutingTableEntryMapper::Map（@0x1fc584e0）、它的两个可达性 worker、公开入口 RoutingTableEntryForICILimitedRouting（@0x1fc58040），以及 GetPhysicalToLogicalMapping3D（@0x1c88a280）均已与函数的 IDA 反编译交叉核对；逐轴 kCaseHopsSignToOffsets 二分查找表和 chip_coordinates 轴字段绑定在下文标为 [LOW] · Part XII — Interconnect & Routing / Routing · 返回索引

摘要

本页记录逐目的地 route-table entry mapper，以及 route table 键控所依赖的物理↔逻辑 core 映射。它负责三个字节精确的机制：

逐 (src,dst) route-table entry — DmaDestinationRoutingTableEntryMapper::Map（@0x1fc584e0），它接受 (source_chip_id, destination_chip_id, ToroidalTopologyInterface&, RoutingScheme)，并返回单个 StatusOr<int> routing-table index：片上路由引擎读取该值来把 DMA 转发到 dst。它根据 RoutingScheme enum 分派到三个可达性 worker 之一（all-to-all direct、n-hop neighbor folding、two-axes）；
physical→logical 3-D map — xla::jellyfish::GetPhysicalToLogicalMapping3D（@0x1c88a280），一个大小为 [Y][X][Z] 的 vector<vector<vector<pair<long,long>>>>，每个叶子为 {core0_logical, core1_logical}；它通过枚举 DeviceAssignment 的每个 (replica, partition) 来填充，先经 TpuTopology::LogicalDeviceForId 把 flat logical-device id 解析为 TpuCoreLocation，再把 logical id 存入该芯片的 (cY,cX,cZ) 坐标。这是 twisted-torus replica-group builder 索引的 device-placement 表；
route-table 数据结构布局 — RoutingScheme enum、把 TpuTopology 包装为 mapper 所消费 ToroidalTopologyInterface 的 slice_builder::Topology adapter，以及 mapper 输出每个 (src,dst) 对一个 index而不是 multi-hop path 的契约。

该 mapper 之上的单播发射层，也就是把这些 entry 串成完整 superpod::routing::RoutingTable 数组的逐 source fiber 扇出（CreateUnicastRoutingTables → CreateSrcDestUnicastRoutingTable → PopulateRoutingTable/GetNextHopAction），不归本页负责；它见 Unicast Route Emission。确定性单路径生成器见 Static-Path Generation；resilient 路径生成器见 Randomized Toroidal Wild-First；预计算路径缓存见 Toroidal Route Cache；逐步 DMA schedule 字面量见 Create-Routing-Schedule。本页从 entry mapper 处接手：它把一个 (src,dst) 转换成一个 index；同时也从该 table 隐式键控的 placement table 处接手。

二进制中观察到的 entry mapper 契约：

Map 返回单个 int routing-table index（chip id），不是 hop list：成功路径写入 *((int*)result + 2) = index 和 *result = 1（ok-Status）。片上路由引擎消费这个 index 来向 dst 转发（见 tpu::RoutingTableEntryForICILimitedRouting 和 net_util::MapSrcDstCoreToRoutingTableIndex）。
mapper 在运行任何 scheme 前都会根据 topology->TotalSize()（芯片数，interface vtable slot +0x78）验证 src 和 dst；越界 chip id 返回 FAILED_PRECONDITION（MakeErrorImpl<9>）。
scheme 是 RoutingScheme enum（a5）：0 = all-to-all (direct)、1 = n-hop（one/two/four/eight neighbors）、2 = two-axes。scheme 0 和 2 带有拓扑大小前置条件（all-to-all ≤ 16 chips；two-axes 仅 2-D 且 ≤ 64 chips，wrap dims 长度为 16）。
GetPhysicalToLogicalMapping3D 是一次纯 placement 复制：每个 (replica, partition) 会解析到其物理芯片坐标，logical id 存储在 mapping[cY][cX][cZ].{first|second}；叶子预初始化为 {-1,-1}（未映射）。per_partition 选择存储的 id 是原始 partition index 还是扁平化的 partition·replicas + replica。

概览

方面	值（byte-anchored）
Entry mapper	`DmaDestinationRoutingTableEntryMapper::Map` `@0x1fc584e0` → `StatusOr<int>`
Mapper 参数	`(int src, int dst, ToroidalTopologyInterface const&, RoutingScheme)`
Scheme enum	`0` = all-to-all · `1` = n-hop · `2` = two-axes
N-hop worker	`MapOneTwoFourEightHopNeighborsReachable` `@0x1fc588a0`
Two-axes worker	`MapTwoAxesReachable` `@0x1fc58fa0`
可达性检查	`CheckReachable` `@0x1fc594c0`
公开入口 (NHop)	`tpu::RoutingTableEntryForICILimitedRouting` `@0x1fc58040`（调用 `Map(...,1)`）
Net-util 调用者 (NHop)	`net_util::MapSrcDstCoreToRoutingTableIndex` `@0x1c6aea80`（调用 `Map(...,1)`）
Topology adapter	`slice_builder::Topology::Topology`（由 `TpuTopology` 在 `@0x1fc58040` 构建）
Interface vtable slot	`+0x48` dim-sizes-cleanup · `+0x50` dim count · `+0x58` per-dim size · `+0x68` is-wrap · `+0x78` TotalSize · `+0x88` GetCoordinate(int) → `StatusOr<Coordinates>`
结果写入 (ok)	`((int)result + 2) = index` · `*result = 1`
Phys→logical map	`xla::jellyfish::GetPhysicalToLogicalMapping3D` `@0x1c88a280`
Map 形状	`vector<vector<vector<pair<long,long>>>>` `[Y=cfg+0x5c][X=cfg+0x58][Z=cfg+0x60]`
Map 叶子	`pair{core0_logical, core1_logical}`，初始化 `{-1,-1}`
Map 消费者	`StrategyND` / `TwistedTorusND` `GetPhase0/1ReplicaGroups`（`@0x137c9e80` / `@0x137cb6a0` / `@0x137cc620` / `@0x137ce240` / `@0x137d3560` / `@0x137d3de0`）
Source TU	`dma_destination_routing_table_entry_mapper.cc` · `group_utils.cc` · `n_hop_route.cc`

1. route-table entry — `DmaDestinationRoutingTableEntryMapper::Map`

1.1 在流水线中的角色

mapper 是 route-table-generation 栈的底部：给定 source chip 和 destination chip，它返回 routing engine 写入 route-table cell 的单个 int，从而让在 src 发起的 DMA 到达 dst。对于 limited ICI routing（n-hop scheme，也是公开入口点请求的唯一 scheme），该 index 是包应该被发送到的下一个芯片，即把 destination chip id 折叠到 torus 的 1/2/4/8-hop neighbor 结构中。单播发射层（Unicast Route Emission）会为完整切片 route table 调用类似的逐 (src,dst) 构建；这个 mapper 是运行时 n_hop_route.cc 路径和 net_util 构建 routing-table-index mapping table 时使用的逐 cell 原语。

Map 是 accel_ssw::deepsea::slice_builder::viperlite_pod::DmaDestinationRoutingTableEntryMapper 上的方法。它的第一个参数（a1）是 StatusOr<int> 的返回槽（absl C++ ABI 通过隐藏指针返回 aggregate）；结果 int 位于 result+0x8（*((int*)result + 2)），Status 判别字位于 result+0x0。

1.2 签名与校验

cpp

// @0x1fc584e0  (dma_destination_routing_table_entry_mapper.cc)
StatusOr<int> DmaDestinationRoutingTableEntryMapper::Map(
    int src,                              // a2
    int dst,                              // a3
    const ToroidalTopologyInterface& topo,// a4 (abstract base; vtable-dispatched)
    RoutingScheme scheme);                // a5  {0,1,2}
```text

前两个操作是根据**芯片数**（`topo.TotalSize()`，interface vtable slot `+0x78`）做边界检查：

```text
n = topo.TotalSize()                        // (*(int(*)(...))(*(qword*)topo + 120))(topo)
if (src < 0 || src >= n)  return INVALID_ARGUMENT("Invalid source chip ID")
if (dst < 0 || dst >= n)  return INVALID_ARGUMENT("Invalid destination chip ID")

两个检查都会调用 interface slot +0x78（+120）取得 n；字符串 "Invalid source chip ID"（len 22，line 34）和 "Invalid destination chip ID"（len 27，line 37）都是 TU dma_destination_routing_table_entry_mapper.cc 中的 MakeErrorImpl<9>（absl kFailedPrecondition = 9，即 FAILED_PRECONDITION，不是 INVALID_ARGUMENT）。（mapper 自身所有拓扑前置条件错误都使用 <9>；只有 not-reachable 和 unsupported-scheme 路径使用 MakeErrorImpl<3> = kInvalidArgument。）

1.3 Scheme 分派 + 前置条件

Map 把 scheme 记录在 v30 中，并在委派前应用 scheme 专属的拓扑前置条件：

`scheme`	名称	前置条件（在 `Map` 中检查）	Worker
`2`	two-axes	dim count == 2（slot `+0x48` 返回 2）且 `TotalSize() <= 64`；每个 is-wrap dim 长度必须为 16（否则 `"All wrap-around dimensions must be of length 16"`）；每条轴必须 ≤ 8（`"Two axes routing must use axes of size <= 8"`）	`MapTwoAxesReachable` `@0x1fc58fa0`
`1`	n-hop	除 src/dst 边界外无其他条件	`MapOneTwoFourEightHopNeighborsReachable` `@0x1fc588a0`
`0`	all-to-all	`TotalSize() <= 16`（否则 `"All to all routing is only supported for slices with <= 16 chips"`）	inline（direct，见 §1.4）
其他	—	—	`INVALID_ARGUMENT("Unsupported routing scheme: %d")`（line 94）

scheme == 2 分支首先检查 (*(...)(*topo + 72))(topo) == 2，即 interface slot +0x48 返回 dim count；错误为 "Two axes routing is only supported for 2-D topologies"（len 53）。随后检查 TotalSize() <= 64，错误为 "Two axes routing is only supported for slices with <= 64 chips"（len 62）。然后它迭代 per-dim size（slot +0x50 total，slot +0x58 per-dim StatusOr<int>），检查 size < 9（"...axes of size <= 8"，line 61），并对 is-wrap dim（slot +0x68 StatusOr<bool>）检查 size >= 16（"All wrap-around dimensions must be of length 16"，line 77）。scheme == 0 && TotalSize() >= 17 guard 先于该循环。尾部分派是 if (v30==2) MapTwoAxes else if (v30==1) MapOneTwoFourEight else if (v30) <error> else <direct>。

1.4 all-to-all direct 情况（`scheme == 0`）

当 scheme 是 all-to-all 且芯片数很小（≤ 16）时，每个芯片都是其他每个芯片的一跳邻居；routing-table index 就是 destination chip id：

text

// scheme == 0, fall-through after the per-dim wrap loop
*((int*)result + 2) = dst;     // routing_table_index = destination chip id
*result = 1;                   // ok Status
```text

> 尾部分派的 `else` 分支（`v30 == 0`）写入 `*((_DWORD*)v32 + 2) = v31`，其中 `v31 = a3 = dst`，并写入 `*v5 = 1`。不会调用 worker，direct routing 表示“直接发送到 destination chip”。

### 1.5 n-hop worker（`MapOneTwoFourEightHopNeighborsReachable`）

这是生产 limited-ICI 路径使用的 worker（每个公开调用者都请求 `scheme == 1`）。它把 `src` 和 `dst` 解析为 torus `Coordinates`（interface slot `+0x88` `GetCoordinate(int) → StatusOr<Coordinates>`），确认 `dst` 可从 `src` 到达（`CheckReachable`，§1.6），然后把 source→destination 位移折叠到 {1,2,4,8}-hop neighbor lattice 来计算**next-hop chip id**。

```text
// @0x1fc588a0
src_coord = topo.GetCoordinate(src)                       // slot +0x88
dst_coord = topo.GetCoordinate(dst)                       // slot +0x88
reach = CheckReachable(src_coord, dst_coord, topo)        // @0x1fc594c0
if (!reach.reachable)
    return INVALID_ARGUMENT("Chip ID %d is not reachable from chip ID %d "
                            "for this topology, %s")        // line 196
// reach carries {reachable: bit32, hops: int}
hops = reach.hops
sx = src_coord.x();  dx = dst_coord.x()
sy = src_coord.y();  axis_len = topo.dimsize(0)            // slot +0x88 with arg 0
GetHopLength(hops)                                         // validates |hops| ∈ {1,2,4,8}
routing_case = (sx == dx) ? ((axis_len <= 4 ? 3 : 1) + (sy&1))   // Y-axis move
                          : ((dx&1) /* X-axis move */ + 1)
sign = (hops <= 0) ? NEGATIVE(2) : POSITIVE(1)
// binary search kCaseHopsSignToOffsets for {routing_case, |hops|, sign}
offsets = kCaseHopsSignToOffsets[{routing_case, hop_len, sign}]  // RET_CHECK if absent
routing_table_index = fold(offsets, sx, sy, axis_len) & 7 | 8    // see [LOW] below
RET_CHECK(routing_table_index != src)                     // line 387
*((int*)result + 2) = routing_table_index
*result = 1

该 worker 通过 slot +0x88 读取 src/dst 坐标，调用 CheckReachable（@0x1fc594c0），并在不可达时用格式字符串 "Chip ID %d is not reachable from chip ID %d for this topology, %s"（len 65，line 196）发出 MakeErrorImpl<3>。可达性结果是 StatusOr<{reachable,hops}>，其中 & 0x100000000 位是 reachable flag，低 dword 是 hop count。会对 hop count 调用 GetHopLength（line 356）。RET_CHECK(routing_table_index != source_chip_id)（字符串 "routing_table_index != source_chip_id"，line 387）防止自环，并带有 VLOG dump（"routing_table_index is equal to source_chip_id for: ..."，line 382）。
[LOW] kCaseHopsSignToOffsets 的精确算术，即在 512 字节区域内对 16 字节 {routing_case, hop_len, sign, offset} 记录执行两层交错二分查找（kCaseHopsSignToOffsets.contains({routing_case, hop_len, hops>0?POSITIVE:NEGATIVE})，line 363），以及最终的 (offset + sy + hops*axis_len) & 7 | 8 折叠（v41 = ((unsigned)offset + v79 + v16*v17 ...) & 7 | 8），已追踪到其操作数，但未化简为逐轴 index 的闭式公式。结构（用 (case, |hops|, sign) 键查表，把存储的 offset 加到带 stride 的坐标上，mask 到该轴，设置 bit 3）已按字节确认；每个 routing_case（0/1/3 + parity）的精确含义来自 sx==dx / axis_len<=4 分支形状推断。

1.6 可达性检查 + two-axes worker

CheckReachable（@0x1fc594c0）接受两个 superpod::routing::Coordinates 和 topology，并返回 StatusOr<{bool reachable, int hops}>，表示 dst 是否是沿单轴从 src 出发的合法 n-hop neighbor，以及带符号 hop count。MapTwoAxesReachable（@0x1fc58fa0）是 scheme == 2 的对应物：它解析两个坐标（slot +0x88），并折叠可能沿两条轴移动的位移（小型 ≤ 64-chip 2-D 情况），同样先校验可达性再生成 index。两个 worker 共享 dma_destination_routing_table_entry_mapper.cc TU 和相同的 StatusOr<int> 返回契约。

MapTwoAxesReachable @0x1fc58fa0 以和 n-hop worker 相同的 GetCoordinate(src)（slot +0x88）模式开头，并且只会从 Map 的 scheme == 2 分支到达。two-axes 前置条件循环（axes ≤ 8，wrap dims == 16）由 Map 在调用前强制执行。

2. 公开入口点

mapper 对 viperlite_pod 私有；运行时代码通过两个 wrapper 到达它，二者都请求 RoutingScheme = 1（n-hop）。

2.1 `tpu::RoutingTableEntryForICILimitedRouting`

cpp

// @0x1fc58040  (n_hop_route.cc — learning/45eac/tpu/runtime/hal/internal/vxc/)
StatusOr<int> RoutingTableEntryForICILimitedRouting(
    const tpu::TpuTopology& topo, int src, int dst);
```text

这是 **HAL 层 limited-ICI route query**。它根据 `topo.chip_count()`（`*((int*)&topo + 28)`）验证 `src`/`dst`，要求**2-D 拓扑**（`*((int*)&topo + 24) == 1`，即 Z dim == 1，否则 `"toplogy must be 2d for limited ICI routing, z: %d"` [sic — 二进制自身拼写如此]），然后从原始 `TpuTopology` 构建 `slice_builder::Topology` adapter 并调用 mapper：

```text
is_wrap_x = topo[+160] & 1
is_wrap_y = topo[+161] & 1
dim_sizes = topo[+88]                       // (qword*)topo + 11
Topology adapter(&dim_sizes, /*ndims=*/2, {is_wrap_x, is_wrap_y}, /*?*/2, ...)
idx = DmaDestinationRoutingTableEntryMapper::Map(adapter, src, dst, /*scheme=*/1)
return ok(idx) ? idx : -1     // on error, returns index -1

该函数用 "Invalid source chip id " / "Invalid destination chip id " 检查 a3 >= topo[+0x70/4]（chip count，int-offset 28 处字段），随后用 "toplogy must be 2d for limited ICI routing, z: " 检查 topo[+0x60/4] != 1（int-offset 24，Z size）。它从 topo[+0xa0]/topo[+0xa1]（wrap bits）和 topo[+0x58]（dim-size span）构造 slice_builder::Topology，然后调用 DmaDestinationRoutingTableEntryMapper::Map(&result, src, dst, &adapter, 1)。成功时写入 *((int*)this + 2) = mapped_index；mapper 返回错误 Status 时写入 -1 并 unref 该错误。

2.2 `net_util::MapSrcDstCoreToRoutingTableIndex`

xla::jellyfish::net_util::MapSrcDstCoreToRoutingTableIndex（@0x1c6aea80）是 compile-side 对应物：它同样构建 topology adapter 并调用 Map(...,1)，供 GenerateRoutingTableIndexMappingTable（@0x1c6a2b80）使用，以物化运行时按 DMA 索引的完整 src×dst routing-table-index table（DMA descriptor 的 routing_table_index 字段；这些 index 如何填充 superpod::routing::RoutingTable 见 Unicast Route Emission）。

MapSrcDstCoreToRoutingTableIndex @0x1c6aea80 包含调用 DmaDestinationRoutingTableEntryMapper::Map(&result, src, dst, adapter, 1)（其反编译第 60 行），确认两个公开调用者都固定 scheme == 1。

3. route-table 数据结构

3.1 mapper 发射什么

mapper 的输出是每个 (src,dst) cell 一个 int32，不是打包的多字段 entry。成功路径统一为：

text

*((int*)result + 2) = routing_table_index;   // the int the routing engine reads
*result = 1;                                  // absl::Status == OK (discriminant 1)
```text

因此运行时编程的完整 route table 是由 `GenerateRoutingTableIndexMappingTable` 对每一对调用 mapper 构建出的密集 `src × dst` `int32` index 数组。更丰富的逐 link `RoutingTable` 行（`unicast_target` / `unicast_terminal` / `vc_control`）由上一层发射层生成；本页的 mapper 提供这些行折入的**next-hop index**。

### 3.2 `ToroidalTopologyInterface` ABI

mapper 从不触碰具体 topology，而是通过抽象 `accel_ssw::deepsea::slice_builder::ToroidalTopologyInterface` vtable 分派，因此任意 slice/pod topology 都能驱动它。`Map` 及其 worker 使用的 slot：

| Vtable 偏移 | 方法（推断） | 返回 | 使用方 |
|---------------|-------------------|---------|---------|
| `+0x28` (`+40`) | topology-name / to-string | `string` | n-hop 错误消息（`%s`） |
| `+0x48` (`+72`) | dim count | `int` | two-axes "must be 2-D" 检查 |
| `+0x50` (`+80`) | dim-sizes accessor (cleanup ptr) | — | per-dim 前置条件循环 |
| `+0x58` (`+88`) | per-dim size | `StatusOr<int>` | axes ≤ 8 / wrap == 16 检查 |
| `+0x68` (`+104`) | is-wrap(dim) | `StatusOr<bool>` | wrap-dim 长度检查 |
| `+0x78` (`+120`) | `TotalSize()` (chip count) | `int` | src/dst 边界，scheme 上限 |
| `+0x88` (`+136`) | `GetCoordinate(int)` | `StatusOr<Coordinates>` | 两个 worker |

> 上述每个 slot 都来自 `Map` / `MapOneTwoFourEightHopNeighborsReachable` 中的 `(*(...)(*(qword*)a4 + N))(a4, ...)` 间接调用。`slice_builder::Topology` 具体类（在 `RoutingTableEntryForICILimitedRouting @0x1fc58040` 构造，vtable `off_220174F0`）是 HAL 路径使用的实现。

### 3.3 `slice_builder::Topology` adapter

`RoutingTableEntryForICILimitedRouting` 将 adapter 构建为 `Topology(dim_size_span, ndims=2, wrap_bits[2], 2, ...)`。adapter 持有两个 dim size、两个 wrap bit（X、Y），并通过 interface vtable 暴露它们。函数退出时销毁该对象（三个 `free` 调用释放 dim-size span、scratch buffer 和 wrap-bit storage）。

---

## 4. physical→logical 3-D map — `GetPhysicalToLogicalMapping3D`

### 4.1 它为什么存在

route table 按**芯片坐标**键控，但 collective 是用**逻辑设备 id**（replica/partition）表达的。`GetPhysicalToLogicalMapping3D` 是反向 placement table，让 twisted-torus replica-group builder 能把物理 `(Y,X,Z)` torus 坐标翻译回占据该芯片的 `{core0, core1}` logical id，也就是 device-assignment 世界与 route-table 世界之间的桥梁。它由 `StrategyND`/`TwistedTorusND` `GetPhase0/1ReplicaGroups`（及其 `NDNway` 变体）消费，这些 builder 会把 twist 坐标折叠为 replica pair。

### 4.2 签名与形状

```cpp
// @0x1c88a280  (group_utils.cc — platforms/xla/service/jellyfish/lowering/)
std::vector<std::vector<std::vector<std::pair<long,long>>>>
GetPhysicalToLogicalMapping3D(const Target& tgt,
                              const DeviceAssignment* da,
                              bool per_partition);

chip_cfg = tgt[+0x3b8]（一个 int*；本页把它当作带有 v5[N] int 字段的 v5）。三层嵌套大小来自 chip config：

层级	大小来源	索引坐标	叶子
outer	`chip_cfg[23]` = `[cfg+0x5c]` = Y	`cY` = `chip_coordinates()` field +4	`vector<vector<pair>>` (24 B)
middle	`chip_cfg[22]` = `[cfg+0x58]` = X	`cX` = `chip_coordinates()` field +0	`vector<pair>` (24 B)
inner	`chip_cfg[24]` = `[cfg+0x60]` = Z	`cZ` = `chip_coordinates()` field +8	`pair<long,long>` (16 B)

每个叶子 pair 初始化为 {-1, -1}：

text

*(qword*)(leaf)      = -1;   // .first  = core0_logical (unmapped)
*(qword*)(leaf + 8)  = -1;   // .second = core1_logical (unmapped)
```text

> outer vector 是 `operator new(24 * Y)`，其中 `Y = v5[23]`；每个 middle vector 通过 `vector<…>::__append` 增长到 `X = v5[22]`；每个 inner 增长到 `Z = v5[24]` 个 `pair`；`{-1,-1}` 初始化是在 `LABEL_34` fill loop 内的两个 `*(qword*)(... ) = -1` store。元素 stride 匹配：outer `24 * v12`，middle `lea [r15+r15*2]`（= `·24`），inner `16 * v77`（该 pair）。

### 4.3 device-assignment 填充循环

```text
replica_count   = da[+0x8] & 0x7fffffff
partition_count = da[+0x0] & 0x7fffffff
for r in [0, replica_count):                 // outer (i)
  for p in [0, partition_count):             // inner (v41)
    // 1. flat logical-device id = Σ coord[i]·stride[i], coord={p, r}, stride=da[+0]
    flat = horner(coord = {p, r}, stride = da[+0x0])     // imul/add chain
    logical_id = da.flat_id_table[flat]      // (int*)(da[+0x10])[flat]
    // 2. resolve to a TpuCoreLocation
    loc = TpuTopology::LogicalDeviceForId(chip_cfg, /*core_type=*/0, logical_id)
    (cY, cX, cZ) = loc.chip_coordinates()    // struct {coord0+0, coord1+4, coord2+8, valid+12}
    cY = field+4;  cX = field+0;  cZ = field+8
    // 3. choose the stored logical id
    stored = per_partition ? p : (p·replica_count + r)
    // 4. choose the megacore slot from the location's "second core" flag
    if (loc.is_second_core /* v75[52] != 0 */)
        mapping[cY][cX][cZ].second = stored      // core1
    else
        mapping[cY][cX][cZ].first  = stored      // core0

循环边界是 da[+0x8] & 0x7fffffff（replica）和 da[+0x0] & 0x7fffffff（partition）。flat-id 是对 {p, r} 坐标和 da[+0x0] stride vector 执行的 Horner 风格 imul/add 链（LABEL_57 处 8-way-unrolled multiply chain / 展开体）。LogicalDeviceForId(chip_cfg, 0, da.flat_id_table[flat]) 使用 core_type = 0 和 da[+0x10] 处的 flat-id table（*((qword*)v38 + 2)）。chip_coordinates() 被读取三次到 v76[1]=cY、v76[0]=cX、v77=cZ。除非 !per_partition 参数，否则 stored id 是 v57 = p + replica_count·r；精确地说，v57 = v88 + v85*r，当 !(byte)v79（per_partition 参数为 false）时 v57 = (unsigned)r。slot 根据 *(dword*)&v75[52]（chip_coordinates 结构体的 valid/second-core 字段）选择：非零 → .second（pair 的 +8），零 → .first。三层都会做边界检查（越界则 BUG()）。
[LOW] chip_coordinates() 背后的精确命名 TpuChipConfig/TpuCoreLocation 字段（coord0=cX、coord1=cY、coord2=cZ，以及 +52 second-core/valid flag）和 per_partition true/false 语义（v57 = p + replicas·r vs v57 = r）在偏移层面已按字节确认；每个偏移到 proto 字段名的绑定来自 VLOG dump 标签推断（group_utils.cc:285 处的 "replica:"、"model id:"、"row:"、"col:"、"dim_z:"）。

4.4 诊断 VLOG

每次填充迭代中，若启用 VLOG(2)（group_utils.cc:285），会发出：

text

device assignment. replica: <r> model id: <p> row: <cY> col: <cX> dim_z: <cZ>
```text

> `LABEL_73` 处的 `LogMessage` 链精确发出这些标签（`"device assignment. replica: "`、`" model id: "`、`" row: "`、`" col: "`、`" dim_z: "`），操作数为 `r`、`p`、`cY`、`cX`、`cZ`，由 `VLogSite::SlowIsEnabled2(..., dword_2236DE58)` 门控。

### 4.5 消费者

该 map 由 replica-group builder 读取，用来把 twist 坐标 `(cY,cX,cZ)` 转换为 `{core0_logical, core1_logical}` 对：

| 消费者 | 地址 |
|----------|---------|
| `StrategyND::GetPhase0ReplicaGroups` | `@0x137c9e80` |
| `StrategyND::GetPhase0ReplicaGroupsNDNway` | `@0x137cb6a0` |
| `StrategyND::GetPhase1ReplicaGroups` | `@0x137cc620` |
| `StrategyND::GetPhase1ReplicaGroupsNDNway` | `@0x137ce240` |
| `TwistedTorusND::GetPhase0ReplicaGroups` | `@0x137d3560` |
| `TwistedTorusND::GetPhase1ReplicaGroups` | `@0x137d3de0` |

> 所有六个函数的反编译体中都引用了 `GetPhysicalToLogicalMapping3D`（通过符号 grep 交叉核对）。这些是 megacore replica-group emitter，会把每个物理芯片的两个 core id 追加到 HLO `ReplicaGroup` device list 中。

---

## 5. `CollectivePermute` route table（交叉引用）

`CollectivePermute` **不会**通过上面的逐 `(src,dst)` mapper 路由。它会直接从 HLO 把 permutation 编译成扁平**source→target id table**（ConstantMapper Type `0xb`）以及逐步 DMA schedule 字面量（Type `5`）：

- `CreateCollectivePermutePairs`（`@0x1347aa40`）读取 `HloInstruction::source_target_pairs()`、`channel_id()` 和 `Target::ReplicaCount()`，生成 `(source_id, target_id)` 的 `vector<pair<long,long>>`。
- `CreateCollectivePermuteTransfers`（`@0x13470fe0`）把每个 id 解码为逻辑坐标，通过 `LogicalTopologyInfo` coordinate→core-id 表（`+0x10`）映射，并为每个 `(pair × buffer × read/write)` 发射一个 16 字节 `net_router::Transfer {src_core@0, src_index@4, dst_core@8, dst_index@0xc}`。
- `CollectivePermuteEmitter::GenerateConstants`（`@0x1346ff60`）把扁平 `s32[]` 表注册为 `AddConstant(0xb, …)`，并把 route schedule（基于同一组 `Transfer`）注册为 `AddConstant(5, …)`。

Type-5 schedule 字面量和 Type-0xb 表布局记录在 collective-lowering 一侧；见 **[Create-Routing-Schedule](create-routing-schedule.md)** 和 **[Collectives Overview](../collectives/overview.md)**。分工如下：Type `0xb` 回答*谁*发给谁（placement permutation，类似上文 `GetPhysicalToLogicalMapping3D` 的 placement keying）；Type `5` 回答*如何/何时*（逐步、逐方向 DMA 程序）。§1 的 route-table-index mapper 是 unicast / auto-routing 路径；CollectivePermute 使用 explicit-schedule 路径。二者最终都会编程 ICI DMA，但机制正交。

---

## 6. 函数映射

| 函数 | 地址 | 作用 |
|----------|---------|------|
| `DmaDestinationRoutingTableEntryMapper::Map` | `0x1fc584e0` | 逐 `(src,dst)` → `StatusOr<int>` routing-table index；scheme 分派 |
| `…::MapOneTwoFourEightHopNeighborsReachable` | `0x1fc588a0` | n-hop（`scheme==1`）worker — 生产 limited-ICI 路径 |
| `…::MapTwoAxesReachable` | `0x1fc58fa0` | two-axes（`scheme==2`）worker（≤ 64-chip 2-D） |
| `…::CheckReachable` | `0x1fc594c0` | `(src_coord, dst_coord)` → 沿轴的 `{reachable, hops}` |
| `tpu::RoutingTableEntryForICILimitedRouting` | `0x1fc58040` | HAL 入口；构建 adapter，调用 `Map(...,1)`，错误时返回 `-1` |
| `net_util::MapSrcDstCoreToRoutingTableIndex` | `0x1c6aea80` | compile-side 入口；调用 `Map(...,1)` |
| `net_util::GenerateRoutingTableIndexMappingTable` | `0x1c6a2b80` | 构建完整 src×dst index table |
| `xla::jellyfish::GetPhysicalToLogicalMapping3D` | `0x1c88a280` | `[Y][X][Z]` → `{core0,core1}` placement map |
| `tpu::TpuTopology::LogicalDeviceForId` | (called) | flat logical id → `TpuCoreLocation`（core_type 0） |
| `tpu::TpuCoreLocation::chip_coordinates` | (called) | `TpuCoreLocation` → `(cX@0, cY@4, cZ@8, valid@12)` |

---

## 7. 诊断字符串

| 字符串（位于 `dma_destination_routing_table_entry_mapper.cc`） | 行 | Status code | 条件 |
|-------------------------------------------------------------|------|-------------|-----------|
| `Invalid source chip ID` | 34 | FAILED_PRECONDITION (`<9>`) | `src` 超出 `[0, TotalSize)` |
| `Invalid destination chip ID` | 37 | FAILED_PRECONDITION (`<9>`) | `dst` 超出 `[0, TotalSize)` |
| `Two axes routing is only supported for 2-D topologies` | 50 | FAILED_PRECONDITION (`<9>`) | `scheme==2`，dim count != 2 |
| `Two axes routing is only supported for slices with <= 64 chips` | 54 | FAILED_PRECONDITION (`<9>`) | `scheme==2`，`TotalSize > 64` |
| `Two axes routing must use axes of size <= 8` | 63 | FAILED_PRECONDITION (`<9>`) | `scheme==2`，axis size ≥ 9 |
| `All to all routing is only supported for slices with <= 16 chips` | 44 | FAILED_PRECONDITION (`<9>`) | `scheme==0`，`TotalSize >= 17` |
| `All wrap-around dimensions must be of length 16` | 77 | FAILED_PRECONDITION (`<9>`) | wrap dim length < 16 |
| `Chip ID %d is not reachable from chip ID %d for this topology, %s` | 196 | INVALID_ARGUMENT (`<3>`) | `CheckReachable` false |
| `Unsupported routing scheme: %d` | 94 | INVALID_ARGUMENT (`<3>`) | scheme 不在 `{0,1,2}` |
| `routing_table_index != source_chip_id` | 387 | RET_CHECK | 计算出的 index == src（自环） |
| `toplogy must be 2d for limited ICI routing, z: %d`（in `n_hop_route.cc`） | 40 | INVALID_ARGUMENT (`InvalidArgumentErrorBuilder`) | Z dim != 1 |

> 上述所有字符串均在两个 TU 中对应的 `MakeErrorImpl` / `InvalidArgumentErrorBuilder` / `RetCheckFailSlowPath` 调用点读取。mapper 自身的六个 precondition error 是 `MakeErrorImpl<9>`（absl `kFailedPrecondition`）；`not-reachable`（line 196）和 `unsupported-scheme`（line 94）路径是 `MakeErrorImpl<3>`（`kInvalidArgument`）；三个 `n_hop_route.cc` 错误（`Invalid source/destination chip id`、`toplogy must be 2d`）是 `util::InvalidArgumentErrorBuilder`。

---

## 交叉引用

- **[Routing Overview](overview.md)** — 该 mapper 所处的 route-generation → cache → emission 流水线底部。
- **[Unicast Route Emission](unicast-route-emission.md)** — 把这些 entry 串成完整 `superpod::routing::RoutingTable` 数组的逐 source fiber 扇出。
- **[Static-Path Generation](get-static-path.md)** — 非缓存发射使用的确定性单路径生成器。
- **[Randomized Toroidal Wild-First](randomized-toroidal-wildfirst.md)** — resilient 路径生成器和 route-cache schema。
- **[Toroidal Route Cache](toroidal-route-cache.md)** — 预计算的逐 `(src,dst)` path cache。
- **[Create-Routing-Schedule](create-routing-schedule.md)** — 逐步 DMA schedule 字面量（Type 5）和 CollectivePermute Type-0xb 表。
- **[Collectives Overview](../collectives/overview.md)** — replica group（通过 `GetPhysicalToLogicalMapping3D` 构建）如何驱动 on-pod collective。

Route-Table 生成 ​

摘要 ​

概览 ​

1. route-table entry — DmaDestinationRoutingTableEntryMapper::Map ​

1.1 在流水线中的角色 ​

1.2 签名与校验 ​

1.3 Scheme 分派 + 前置条件 ​

1.4 all-to-all direct 情况（scheme == 0） ​

1.6 可达性检查 + two-axes worker ​

2. 公开入口点 ​

2.1 tpu::RoutingTableEntryForICILimitedRouting ​

2.2 net_util::MapSrcDstCoreToRoutingTableIndex ​

3. route-table 数据结构 ​

3.1 mapper 发射什么 ​

4.4 诊断 VLOG ​

Route-Table 生成

摘要

概览

1. route-table entry — `DmaDestinationRoutingTableEntryMapper::Map`

1.1 在流水线中的角色

1.2 签名与校验

1.3 Scheme 分派 + 前置条件

1.4 all-to-all direct 情况（`scheme == 0`）

1.6 可达性检查 + two-axes worker

2. 公开入口点

2.1 `tpu::RoutingTableEntryForICILimitedRouting`

2.2 `net_util::MapSrcDstCoreToRoutingTableIndex`

3. route-table 数据结构

3.1 mapper 发射什么

4.4 诊断 VLOG