AllToAll 表

此页面上的所有地址适用于 libtpu-0.0.40-cp314 轮中的 libtpu.so(内部版本号 89edbbe81c5b328a958fe628a9f2207d，内部版本 libtpu_lts_20260413_b_RC00)。其他版本会有所不同。 .text VMA等于文件偏移量；下面的每个地址都是 VMA。

摘要

xla::jellyfish::GenerateAllToAllTables (0x133ed620) 是 ICI all-to-all 集体的 障碍参与者表 的编译时构建器。它不构建每步合作伙伴计划——即由路由子系统 (创建路由计划) 生成的 Type-5 路由文字。相反，它回答了副本组内屏障所需的正交问题：谁在我的组中，我在其中坐在哪里？ 该函数返回一个 tuple<vector<int> A, vector<int> B, optional<vector<int>> C>，AllToAllEmitterBase::GenerateConstants (0x10f089a0) 将其注册到每个指令 ConstantMapper 中，作为 ConstantMapper 类型 8、9 和0xa，而 GetConstantTables (0x10f07860) 稍后回读并交给 BarrierWithinReplicaGroupStartNoReturn 作为 (InfoTable A, InfoTable B, optional<InfoTable> C)。

了解 MPI 的读者拥有该框架：在环/递归集体中，每个等级必须知道其通信器成员身份(与它同步的对等点集合及其在该集合中的序号)，独立于每次迭代的发送/接收计划。表 A 和 B 正是该成员资格映射及其转置：A 是 device_id → (replica_group, within-group ordinal)(每个设备 2 个整数)，B 是转置枚举 (within-group position, group) → device_id。表 C 是可选的 3-D/分区感知核心表，仅存在于也具有静态设备分配的通道 ID 奇数路径上。相同的三个表驱动 ragged-all-to-all 变体，它共享成员资格生成器和全对所有成本帮助器，仅在运行时偏移算术方面有所不同。

本页记录了三个可重新实现的工件：(1) GenerateAllToAllTables 中的表生成循环(组大小选择、副本组填充和无副本组标识填充、表 C 跨步线性化、返回元组布局)； (2) 通过GenerateConstants→ConstantMapper→GetConstantTables的注册/回读路径以及静态与动态载体选择； (3) 每个 XY 路由缓冲区 Allocator 值结构 — 路由调度回调维护的 FlatHashMap<XY, Allocator> 暂存缓冲区记分板，其值布局 (int32 size + std::deque<pair<int,int>> available) 在这里按字节精确解码。网格几何、成员资格表索引 (MeshNDInfo) 和共享该结构的 ND 环 AllGather 副本表在它们接触所有内容的地方内联记录，并由 AllGather ND 环完全拥有。

重新实现，合约为：

成员编码。 表 A(Type 8，2 个整数/设备)、B(Type 9，1 个整数/槽)和可选 C(Type 0xa)；它们的索引算法、它们的存储值以及两个填充路径(副本组与身份)。
组尺寸和 C 桌门。 group_size = channel_id & 1 ? mesh_dim1 : mesh_dim0； static_device_assignment 前提条件； -0x68“表 C 存在”标志仅在通道 ID 奇数 + 副本计数路径上设置。
运营商选择。 CollectiveShouldUseStaticInfoTable (0x138194c0) 根据 TpuCompEnv[+0x15d0] >= mesh_dim0·mesh_dim1 阈值选择每个表的静态 R1 Literal 与运行时 vector<int>。
分配器值结构。 FlatHashMap<XY, Allocator> 0x40 字节槽(8 字节 XY 键 + 0x38 字节值)、四个释放时间 RET_CHECK 不变量和单个 available 双端队列(没有单独的“最新 DMA 输出”容器)。


建表器	`xla::jellyfish::GenerateAllToAllTables(const HloInstruction&, long, long)` @ `0x133ed620`
SparseCore 双	`GenerateAllToAllTablesForSparseCore(const HloInstruction&, const DeviceAssignment&)` @ `0x133ee200`
注册商	`AllToAllEmitterBase::GenerateConstants` @ `0x10f089a0`(类型 8 / 9 / 0xa + 类型 5)
运输门	`CollectiveShouldUseStaticInfoTable` @ `0x138194c0` (`TpuCompEnv[+0x15d0] >= mesh_dim0·mesh_dim1`、`setge`)
回读	`(anon)::GetConstantTables` @ `0x10f07860` → `(InfoTable, InfoTable, optional<InfoTable>)`
屏障消费者	`AllToAllEmitterBase::EmitBarrierStartImpl` @ `0x10f07240` → `BarrierWithinReplicaGroupStartNoReturn` @ `0x1c6983e0`
成本分支	`ComputeAllToAllCycles` @ `0x130ae8e0` / `ComputeRaggedAllToAllCycles` @ `0x130aea80` (共享 `ComputeAllToAllCyclesHelper` @ `0x130d02a0`)
路由缓冲区记分板	`FlatHashMap<XY, Allocator>` 通过 `find_or_prepare_insert` @ `0x138270a0`(槽步长 `0x40`)； `$_1` 释放回调@`0x13826dc0`
来源 TU	`platforms/xla/service/jellyfish/lowering/all_to_all_emitter_*.cc`(str@`0x878b4af`)； `net_router_emitter.cc` 中的分配器(str @ `0x8760f44`)
信心	HIGH — 构建器、SparseCore twin、读回和分配器回调的反编译验证主体；低行标记为内联

这在哪里

GenerateAllToAllTables 是全对全集合编译成的两个正交常数族之一。拆分是此页面的中心结构事实：

text

HLO all-to-all (opcode 12) / ragged-all-to-all (opcode 86)
        │
        ├── PARTNER SCHEDULE  (the per-(core, step, direction) DMA program)
        │     CreateAllToAllTransfers → CreateRoutingScheduleLiteral
        │       → ConstantMapper Type 5
        │       → CalculateWithLimitedIciRouting / EmitForLimitedIciRouting
        │     (owned by:  ../routing/create-routing-schedule.md)
        │
        └── BARRIER MEMBERSHIP  (who is in my replica group, and my ordinal)
              GenerateAllToAllTables  →  Types 8 / 9 / 0xa
                → GetConstantTables
                → BarrierWithinReplicaGroupStartNoReturn
              (THIS PAGE)
```text

两者由单独的构建器生成并由单独的代码使用。 Type-5 路由文字回答*每步哪个核心 DMA 给谁*；类型 8/9/0xa 答案*谁与谁障碍* — 成员资格集 `GetReplicaGroupCoreInfo` 读取以恢复每个核心的副本组对等集和主控(主控 = 对等 0)。全部 **成本**(`ComputeAllToAllCycles` @ `0x130ae8e0`，除以 `EstimatePhysicalLinksUsed`)是第三个正交轴 - 请参阅 [SPMD 链路计数成本](spmd-link-count-cost.md)。

该页面拥有**会员表生成器、其注册/读回以及路由缓冲区 `Allocator` 记分板结构**。 AllGather ND 环副本表(`CreateStaticNDRingReplicaInfoTable`，ConstantMapper 类型 0/1/2)共享 `MeshNDInfo` 几何结构和相同的 `ConstantMapper` 机制，但属于不同的集合；它们归 [AllGather ND 环](allgather-nd-ring.md) 所有。甚至到达该构建者的策略决策都在[SelectND策略](strategy-nd-picker.md)上。

---

## 1. `GenerateAllToAllTables` — 表生成循环

### 用途

为一个全对所有 `HloInstruction` 构建三个成员资格表。该函数采用指令加上两个网格范围(`mesh_dim0`、`mesh_dim1`)，并通过 `sret` 返回 `(A, B, optional C)` 元组。它是 `CreateStaticReplicaInfoTable` 的全面模拟：它在编译时预先计算运行时屏障必须导出的设备到组映射。

### 入口点

```text
AllToAllEmitterBase::GenerateConstants @0x10f089a0
  └─ GenerateAllToAllTables @0x133ed620                    ── builds (A, B, opt C)
       ├─ HloInstruction::channel_id                       ── parity selects group_size
       ├─ HloInstruction::has_replica_groups / replica_groups
       ├─ GetModule()->config[+0x20][+0x660]               ── static_device_assignment has-bit
       └─ proto2::LogIndexOutOfBoundsAndAbort @0x21063300  ── bounds trap on table writes

算法

两个网格范围通过 AllToAllEmitterBase::GenerateConstants 到达，AllToAllEmitterBase::GenerateConstants 从 LogicalTopologyInfo(mesh_dim0 = movslq (r15)、mesh_dim1 = movslq 0x4(r15))读取它们并调用 GenerateAllToAllTables(hlo, mesh_dim1, mesh_dim0)。在构建器内部，group_size 是 channel-id parity 选择：

function GenerateAllToAllTables(hlo, arg_dim1, arg_dim0):     // 0x133ed620, sret -> (A, B, opt C)
    // group_size = odd channel -> mesh_dim1, even -> mesh_dim0   (cmovne @0x133ed669)
    group_size = (channel_id(hlo) & 1) ? arg_dim1 : arg_dim0;  // line 94 / 121-125

    // The 3-D / partition-aware path requires a static device assignment.
    has_static_da = module_config[+0x20][+0x660];              // movzbl @0x133ed67d
    if ((channel_id(hlo) & 1) && replica_count_present && !has_static_da)
        LogMessageFatal("static_device_assignment_.has_value()");  // hlo_module_config.h:285 @0x133ee1c0

    table_C_present = (channel_id(hlo) & 1) && replica_count_present;  // -0x68 set @0x133edc42

    // ---- Table A : Type 8 ----  vector<int>, 2 ints per device id
    table_A.reserve(2 * group_size);                           // __append @line 168

    if (!has_replica_groups(hlo)) {
        // identity / single-group form
        for (pos in [0, group_size)):
            table_A[2*pos]   = 0;                              // group  = 0   @0x133edbd8
            table_A[2*pos+1] = pos;                            // ordinal = pos @0x133edbf3
            table_B[pos]     = pos;                            // identity      @0x133edc0b
    } else {
        // replica-group form
        for (g, group) in enumerate(replica_groups(hlo)):     // members inline@0x18 / heap@0x20, count@0x1c
            for (pos, dev) in enumerate(group.replica_ids):
                table_A[2*dev]   = g;                          // @0x133edd36
                table_A[2*dev+1] = pos;                        // @0x133edda1
                table_B[group_size*pos + g] = dev;             // imul %r8d,%eax; +pos @0x133edde5
    }

    // ---- Table C : Type 0xa (optional) ----  built only when table_C_present
    if (table_C_present):
        for (pos in within-group positions):
            coord     = (...);                                 // within-group position
            device_id = dot(coord, DA_per_dim_strides);        // 8-wide imul/add @0x133eda50/daf0/df50
            table_C[pos] = device_id;                          // copied from -0xb8 snapshot @0x133ee011

    return { tag=1, A, B, C, has_value = table_C_present };    // sret tuple @0x133ee03a
```text

重新实现者必须正确处理的三个细节：

- **表 A 为 2 宽。** 每个设备 ID 消耗两个连续的 `int` 插槽：`table_A[2·dev]` = 副本组，`table_A[2·dev+1]` = 组内序号。商店站点 `0x133edd36` / `0x133edda1` 和 168 号线的 `__append(2 * group_size)` 预订确认了这一进展；边界检查 `*(_QWORD*)&v83[8] <= 2*v27`(第 263 行)是每对写入之前的 `vector::size()` 防护。
- **表 B 是转置。** 它由 `group_size·position + group`(`0x133edde5` 处的 `imul %r8d(group_size),%eax; add pos`)索引并存储设备 id。读取 A 给出“设备 *d* 位于何处”；读数 B 给出“哪个设备位于(位置，组)”。屏障使用 B 按顺序枚举对等点，使用 A 查找自己的序数。
- **表 C 是有条件的和 3-D。** 仅当通道 id 为奇数 *并且* 存在副本计数(`-0x68` 设置为 `0x133edc42`)时，它才存在，这正是断言 `static_device_assignment_.has_value()` 的路径。它的值是由 8 宽跨步点积与 `DeviceAssignment` 每维跨步数组生成的平面设备 ID — 与 ND 环构建器使用的坐标线性化相同(第 4 节)。

> **明白了 —** 这些是*会员资格*表，而不是合作伙伴计划。将表 B 视为“我在步骤 *k* 发送到的设备”的重新实现者将构建一个损坏的集合。 B为静态枚举`(position, group) → device`；每步发送/接收伙伴来自 [创建路由计划](../routing/create-routing-schedule.md) 中的 Type-5 路由文字。这两个表由不同的函数生成，并提供不同的运行时入口点(屏障与有限 ICI 路由)。

### 返回元组布局

`-0xd0` 处的 `sret` 缓冲区是 `tuple<vector<int>, vector<int>, optional<vector<int>>>`：

| 字段 | 偏移 | 型 | 含义 |
|---|---|---|---|
| 标签 | `+0x00` | int | 活跃成员标签(此处始终为 `1`) |
| A | `+0x08` / `+0x10` / `+0x18` | `vector<int>` `{data, size, cap}` | 表 A(设备→组，序号) |
| B | `+0x20` / `+0x28` / `+0x30` | `vector<int>` `{data, size, cap}` | 表B(位置、组→设备) |
| C | `+0x38` / `+0x40` / `+0x48` | `vector<int>` `{data, size, cap}` | 表C(3维核心表) |
| C.has_value | `+0x50` | 字节 | 仅当 `-0x68` 时设置 `1` @ `0x133ee0ec` |

### 功能图

| 功能 | 地址 | 角色 |
|---|---|---|
| `GenerateAllToAllTables` | `0x133ed620` | 构建`(A, B, opt C)`； 622行反编译行 |
| `GenerateAllToAllTablesForSparseCore` | `0x133ee200` | SparseCore twin(无奇偶校验，无 C)；确认 A/B 编码 |
| `proto2::LogIndexOutOfBoundsAndAbort` | `0x21063300` | 表写入上的边界陷阱(`0x133edd71`、`0x133ee01e`) |
| `LogMessageFatal`(静态 DA 断言) | — | `hlo_module_config.h:285` @ `0x133ee1c0`，str `0x86508f3` |

### 注意事项

A 和 B 的“哪个轴是副本，哪个是分区”标签是**属性，而不是自描述**(低)：索引算术和存储值是字节确认的，但符号含义(A = 副本轴序数表，B = 分区轴枚举)是从 `GetReplicaGroupCoreInfo` 在屏障处读取它们的方式恢复的，而不是从建设者。将*编码*视为高，将*轴标签*视为结构读数。

---

## 2. SparseCore双胞胎——独立A/B确认

### 用途

`GenerateAllToAllTablesForSparseCore` (`0x133ee200`) 是构建器的更简单的 SparseCore 变体。它**没有通道 ID 奇偶校验**和 **没有表 C** - 它只生成 A 和 B - 并且它从集体操作组模式而不是从原始 `replica_groups()` 派生其参与者。它在这里的价值是佐证：它从完全不同的控制路径重现完全相同的 A/B 索引算术，这修复了单个构建器之外的编码。

### 算法

```c
function GenerateAllToAllTablesForSparseCore(hlo, da):        // 0x133ee200
    total = da[+0] * da[+8];                                  // mesh extents product @0x133ee234
    mode  = GetCollectiveOpGroupMode(hlo);                    // 0x1e46bac0, line 67
    groups = GetParticipatingDevicesGroups(da, ..., mode);    // 0x1e46bc20, line 77
    for (g, group) in enumerate(groups):
        group_size = group.size();
        for (pos, dev) in enumerate(group):
            table_A[2*dev]   = g;                             // @0x133ee35f, line  (2*v21)
            table_A[2*dev+1] = pos;                           // @0x133ee37d
            table_B[group_size*g + pos] = dev;                // @0x133ee39d

商店站点(0x133ee35f / 0x133ee37d / 0x133ee39d)和相同的 *(_QWORD*)&v37[8] <= 2*v21 边界保护(第 100 行)与主构建器的编码完全匹配。与§1唯一的结构差异是SparseCore孪生索引B为group_size·group + pos，而密集构建器使用group_size·pos + group——转置约定翻转，因为SparseCore循环嵌套组外层而不是位置外层；两者都将设备 ID 存储在相应读取器期望的插槽中。

注意 — SparseCore 孪生使得 A/B 编码具有高置信度而不是单一见证。两个独立的构建器，两个不同的参与者派生路径，相同的 2·dev / 2·dev+1 商店。

3. 注册和回读 — 类型 8 / 9 / 0xa

用途

AllToAllEmitterBase::GenerateConstants 将三个向量注册到固定类型标签下的每条指令 ConstantMapper 中，选择每个表是否存储静态 R1 Literal(运行时直接读取的烘焙常量)或动态 vector<int>(在运行时实现)。当 InfoTable 和 EmitBarrierStartImpl 将它们反馈到副本组内屏障时，GetConstantTables 会读回它们。

算法

function AllToAllEmitterBase::GenerateConstants(hlo, target, topo, region):  // 0x10f089a0
    mesh_dim0 = movslq (topo);  mesh_dim1 = movslq 0x4(topo);   // 0x10f08a02 / a0a
    (A, B, C) = GenerateAllToAllTables(hlo, mesh_dim1, mesh_dim0);  // 0x10f08a10

    // static carrier iff TpuCompEnv[+0x15d0] >= mesh_dim0 * mesh_dim1
    use_static = CollectiveShouldUseStaticInfoTable(topo);      // 0x138194c0, setge

    AddConstant(Type=8, use_static ? Literal(A) : vector(A));   // 0x10f08cc8 / 0x10f08c2c
    AddConstant(Type=9, use_static ? Literal(B) : vector(B));   // 0x10f08d5b / 0x10f08e27
    if (C.has_value)
        AddConstant(Type=0xa, use_static ? Literal(C) : vector(C));  // 0x10f08ee4 / 0x10f08fa2

    // the orthogonal partner schedule:
    AddConstant(Type=5, CreateAllToAllRoutingScheduleTable(...));   // AddConstant @0x10f09161; call @0x10f0906f -> 0x10f061c0
```text

`CollectiveShouldUseStaticInfoTable`(`0x138194c0`)计算`total = movslq(rsi) · movslq 0x4(rsi)`(= `mesh_dim0·mesh_dim1`)，读取`GetTpuCompEnv`，并通过`cmp %rbx,0x15d0(%rax); setge`返回`TpuCompEnv[+0x15d0] >= total`。低于阈值时，表足够小，可以作为常量进行烘焙；在它之上，它们被动态地具体化。

读回关闭循环：

```c
function GetConstantTables(hlo, mapper):                       // 0x10f07860
    A = mapper.GetConstant(Type=8);                            // 0x10f07889
    B = mapper.GetConstant(Type=9);                            // 0x10f078d2
    optional C;
    if (mapper.HasConstant(Type=0xa))                          // 0x10f07907
        C = mapper.GetConstant(Type=0xa);                      // 0x10f07920
    return StatusOr<tuple<InfoTable, InfoTable, optional<InfoTable>>>(A, B, C);  // 0x10f079f1..a34

function EmitBarrierStartImpl(...):                            // 0x10f07240
    (A, B, C) = GetConstantTables(hlo, mapper);                // 0x10f074e5
    BarrierWithinReplicaGroupStartNoReturn(                    // 0x1c6983e0, call @0x10f07635
        ..., /*rcx*/ A, /*r8*/ B, /*r9*/ C);                   // A=-0xd0, B=-0xb8, C=-0x60

HasConstant(Type=0xa) 在反编译的 GetConstantTables (HasConstant(a2, 10)) 中进行字节确认，因此读取侧门表 C 的存在与构建器门控其构造完全相同。三个 InfoTable 直接位于屏障启动参数寄存器 rcx/r8/r9 中。

功能图

功能	地址	角色
`AllToAllEmitterBase::GenerateConstants`	`0x10f089a0`	寄存器类型 8/9/0xa + 类型 5
`CollectiveShouldUseStaticInfoTable`	`0x138194c0`	静态与动态载流门
`(anon)::GetConstantTables`	`0x10f07860`	将 8/9/0xa 读回为 `InfoTable` 三重
`AllToAllEmitterBase::EmitBarrierStartImpl`	`0x10f07240`	提供三倍至障碍启动
`BarrierWithinReplicaGroupStartNoReturn`	`0x1c6983e0`	集团内障碍消费者
`CreateAllToAllRoutingScheduleTable`	`0x10f061c0`	正交 Type-5 路由文字(调用站点 `0x10f0906f`)

QUIRK — 载体选择是按指令进行的，而不是按表类型进行的。所有三个 A/B/C 从一个 CollectiveShouldUseStaticInfoTable 调用中采取相同的 use_static 决策，因此一条指令永远不会将静态 A 与动态 B 混合。阈值 TpuCompEnv[+0x15d0] 是一个编译环境旋钮，因此相同的 HLO 可以在一个配置上降低到静态常量，在另一个配置上降低到运行时向量。

4. `MeshNDInfo`几何表索引

表 C 的 3-D 线性化和 ND 环 AllGather 表共享一个网格几何描述符。 MeshNDInfo(复制品0x127b5100)是0x40字节，描述嵌入在N维设备网格中的一个每轴环：

字段	偏移	型	含义
轴 ID	`+0x00`	`vector<MeshDim>` (int32×)	网格轴id列表； `memcpy 4·n` @ `0x127b516b`
每个尺寸尺寸	`+0x18`	`vector<long>`(8字节×)	环长度；模邻居的除数； `memcpy 8·n` @ `0x127b51a9`
环令	`+0x28`	`vector<MeshDim>` (int32×)	遍历顺序(沿环的设备id)； `memcpy 4·n` @ `0x127b51e9`
暗淡位掩码	`+0x38`	长	`popcount(low 3 bits)` ⇒ `Is2D` (2) / `Is3D` (3)

表 C (§1) 和 ND 环生成器使用的坐标到设备线性化是网格坐标与 DeviceAssignment 每维度步幅数组的 8 宽 imul/add 点积。在 CreateStaticNDRingReplicaInfoTable (0x1c69e900) 中，调度由描述符为 2-D 或 3-D 的 RET_CHECK 门控：__popcnt(a2[7] & 7) != 2 然后 ... || mesh_info.Is3D()，在失败时报告 net_util.cc:2440 (str 0xa17c039)。 DA 展平边界是 indexes.size() == num_dimensions()(str 0xa1567ac，反编译的第 413 行)。

注意 — MeshNDInfo 和 ND 环构建器(CreateStaticNDRingReplicaInfoTable @ 0x1c69e900、CreateNDRingReplicaInfoTable @ 0x1c69e7e0，后者将前者包裹在 LiteralUtil::CreateR1<int> 的 net_util.cc:2412 中)属于 AllGather ND 环路径(ConstantMapper 类型 0/1/2)。它们出现在这里只是因为表 C 重复使用了相同的 DeviceAssignment 跨步线性化。它们的完整每轴环邻编码和 AllGatherEmitter::InitDim 轴行走归 AllGather ND 环所有。

5. 路由缓冲区 `Allocator` 值结构

用途

路由调度求解器在执行传输调度时维护临时缓冲区的每个目的地记分板。该容器是FlatHashMap<net_util::XY, net_router::(anon)::Allocator>；它的值结构跟踪目标 XY 可以使用的缓冲区数量以及当前正在运行的缓冲区数量。此页面从读取它的缓冲区释放回调中解码值布局字节精确，解决了之前的开放问题“可用”和“最新DMA输出”是否是两个容器(它们是一个)。

插槽和值布局

find_or_prepare_insert (0x138270a0) 使用 0x40 字节槽(步长 shl $0x6 @ 0x13827163 / 0x13827217)。该槽是pair<const XY, Allocator>：slot+0x0处的8字节XY密钥(2×int32)，Allocator值在slot+0x8处。

价值关闭	插槽关闭	字段	型	角色
`+0x00`	`+0x08`	`size`	int32	每 XY 暂存缓冲区计数范围 (`RET_CHECK *ptr.index < size`)
`+0x08`	`+0x10`	`available`	`std::deque<pair<int32,int32>>`	飞行中 `(buffer_index, available_at_step)` 列表

双端队列(libc++ 分割缓冲区，0x30 字节)在值内布局为 __map_.__first_ @ value+0x8、__begin_ @ +0x10、__end_ @ +0x18、__end_cap_ @ +0x20、__start_ @ +0x28、__size_ @ +0x30。 512 元素块索引(shr $0x9、and $0x1ff)是 std::deque 映射签名。总值 = 0x38 字节(int32 size + 0x30 双端队列)，恰好填充 8 字节 XY 密钥后的 0x40 槽。

释放回调不变量

$_1 缓冲区释放回调 (0x13826dc0) 是最终读取器。在每个版本中，它都会查找或插入目标 XY(通过 vmovups ymm0 对插入时的值进行零初始化)，计算 available_at = step + 1，然后在将释放的缓冲区推送到 available 之前强制执行四个 RET_CHECK：

function release_buffer(map, ptr, dest_xy, step):              // $_1 @0x13826dc0
    alloc = map.find_or_insert(dest_xy);                       // fopi @0x13826de1
    available_at = step + 1;                                   // inc %r15d @0x13826e14/29

    RET_CHECK(alloc.available.empty() ||                       // net_router_emitter.cc:389
              alloc.available.back().second <= available_at);  //   sorted by release step @0x13826e50
    RET_CHECK(ptr.type == PointerType::kAlloc);                // :390  cmp $2 @0x13826e8d
    RET_CHECK(ptr.index.has_value());                          // :391  test $1 @0x13826e97
    RET_CHECK(c_none_of(alloc.available,                       // :395  deque scan @0x13826ea4..f1c
                        [&](pair p){ return p.first == *ptr.index; }));  // no double release
    RET_CHECK(*ptr.index < alloc.size);                        // :396  cmp %r14d,(%rax) @0x13826f1e

    alloc.available.push_back( *ptr.index | (step << 32) );    // __add_back_capacity @0x13826f4d
                                                               //   64-bit store @0x13826f76
```text

所有五个字符串在反编译时均进行字节确认： `net_router_emitter.cc:389` `"available.empty() || available.back().second <= available_at"`、`:390` `"ptr.type == PointerType::kAlloc"`、`:391` `"ptr.index.has_value()"`、`:395` `c_none_of` 拉姆达，`:396` `"*ptr.index < size"`。推送将释放的缓冲区打包为 `(index | step<<32)` — `.first` = 缓冲区索引，`.second` = 释放步骤 — 保留 back() 按步骤排序的不变量。

路由缓冲区分配器在 `value+0x8` 保留单个 `std::deque<pair<int,int>> available` — 没有第二个容器；双端队列的`.second`字段携带释放步骤。

### 注意事项

`size` 字段的 **writer** 不是隔离的(低)：它是字节确认的 *读取* 作为 `*ptr.index < size` 检查的每个 XY 绑定，并在第一次插入时初始化为零，但此处未跟踪将其设置为拓扑派生的暂存缓冲区计数的路由调度主循环站点。 `$_1` 捕获是 `0x28` 字节 POD(@ `0x13820aa1` 构建)：`{ &FlatHashMap @+0x00; ptr{type,index} 16B @+0x08; dest XY @+0x18; step int32 @+0x20 }`。拥有此记分板的完整路由调度驱动程序位于 [创建路由计划](../routing/create-routing-schedule.md) 上。

---

## 6. 衣衫褴褛的变种

`ragged-all-to-all`(HLO 操作码 86)是一个依赖于数据的 all-to-all，其每列发送/接收计数是运行时值而不是编译时统一分片。对于 **membership** 表，它与普通的 all-to-all 相同：它是一个 all-to-all-family 操作，因此障碍参与者集来自相同的 `GenerateAllToAllTables` 构建器(类型 8/9/0xa)和相同的 `GetConstantTables` 读回。不同之处在于数据路径上的“运行时偏移算术”(索引到每个对等方缓冲区的参差不齐的偏移量)，由路由调度和 DMA 发射处理，而不是此页面上的成员资格表。

成本模型将两者视为一个系列：`ComputeRaggedAllToAllCycles` (`0x130aea80`) 和 `ComputeAllToAllCycles` (`0x130ae8e0`) 都称为共享的 `ComputeAllToAllCyclesHelper` (`0x130d02a0`)，并且都除以`EstimatePhysicalLinksUsed` (`0x1c8939c0`) — 全链路饱和成本形状。每种成本的公式位于 [SPMD 链路计数成本](spmd-link-count-cost.md) 上。

> **注意 —** 成员资格表构建器没有对 ragged 进行特殊处理；分歧完全在于运行时偏移路径。重新实现者可以对两个操作码使用一个构建器，并且仅在路由调度/DMA 偏移发射中使用分支。

---

## 验证说明

> 表构建器、SparseCore twin、注册/读回和分配器值结构与 `libtpu.so` v0.0.40 的 IDA 反编译进行了交叉检查：
>
> - `GenerateAllToAllTables` @ `0x133ed620`(622行)：`channel_id`奇偶校验(第94/121-125行)，`static_device_assignment_.has_value()`断言(`hlo_module_config.h:285`)， `has_replica_groups`/`replica_groups()`、`__append(2*group_size)`(第 168 行)、表 A `2·dev` 存储(第 263-266 行、`0x133edd36`/`da1`)、表 B `imul group_size` (`0x133edde5`)、表 C 副本 (`0x133ee011`)、`LogIndexOutOfBoundsAndAbort` (`0x21063300`) — 全部存在。
> - `GenerateAllToAllTablesForSparseCore` @ `0x133ee200`：`GetCollectiveOpGroupMode`(第 67 行)、`GetParticipatingDevicesGroups`(第 77 行)、`2·dev`/`2·dev+1` 存储(第 100-103 行)— 准确。
> - `GetConstantTables` @ `0x10f07860`：两个 `GetConstant` 读取 + `HasConstant(a2, 10)` 门，用于 0xa 类型 — 准确。
> - `MeshNDInfo` copy ctor @ `0x127b5100`：三个向量(`4·n`，`8·n`，`4·n` memcpy)+暗长； `CreateStaticNDRingReplicaInfoTable` @ `0x1c69e900` `Is2D/Is3D` popcount (`net_util.cc:2440`), `indexes.size() == num_dimensions()`; `CreateNDRingReplicaInfoTable` @ `0x1c69e7e0` `CreateR1<int>` 在 `net_util.cc:2412` — 准确。
> - `$_1` 释放回调 @ `0x13826dc0`：`RET_CHECK` 第 389/390/391/395/396 行及其完整字符串 (`net_router_emitter.cc`) — 准确；经过验证的单个 `available` 双端队列(无第二个容器)。
> - 成本：`ComputeAllToAllCycles` @ `0x130ae8e0` 和 `ComputeRaggedAllToAllCycles` @ `0x130aea80` 都调用 `ComputeAllToAllCyclesHelper` @ `0x130d02a0` — 准确。
>
> **[LOW]** (1) A/B 轴标签(A = 副本轴，B = 分区轴)归因于屏障的 `GetReplicaGroupCoreInfo` 读取模式，而不是命名字段。 (2)`Allocator.size`写入者未被隔离；仅确认其读取语义和零初始化。 (3) `MeshNDInfo +0x38` 位→轴映射为结构读取；只有 popcount→dimension-count 是字节确认的。

---

## 相关组件

| 名称 | 关系 |
|---|---|
| `CreateAllToAllRoutingScheduleTable` @ `0x10f061c0` | 正交 Type-5 合作伙伴时间表(路线，非会员资格) |
| `BarrierWithinReplicaGroupStartNoReturn` @ `0x1c6983e0` | `(A, B, opt C)` `InfoTable` 三重的消费者 |
| `CreateStaticNDRingReplicaInfoTable` @ `0x1c69e900` | AllGather ND 环表共享 `MeshNDInfo` / DA 跨步线性化 |
| `ComputeAllToAllCyclesHelper` @ `0x130d02a0` | 共享所有操作码的成本助手 |

## 交叉引用

- [On-Pod 集体 — 剖面图](overview.md) — 集体堆栈映射；全对全 (12) / 参差不齐 (86) 成本分支
- [AllGather ND 环](allgather-nd-ring.md) — `MeshNDInfo`，ND 环复制工作台(类型 0/1/2)，`InitDim` 轴行走
- [SelectND策略](strategy-nd-picker.md) — 决定此构建器上游 ICI 环算法的选择器
- [SPMD 链路计数成本](spmd-link-count-cost.md) — `ComputeAllToAllCycles` / `ComputeRaggedAllToAllCycles`、`EstimatePhysicalLinksUsed` 除数
- [常量映射器](constant-mapper.md) — 存储类型 5/8/9/0xa 的每条指令 `ConstantMapper`
- [创建路由计划](../routing/create-routing-schedule.md) — Type-5 路由文字和路由缓冲区 `Allocator` 驱动程序
- [路由表生成](../routing/route-table-generation.md) — 每颜色 `RingLocation` 邻居时间表
- [屏障](../barrier/overview.md) — `BarrierWithinReplicaGroup`、`GetReplicaGroupCoreInfo`，成员资格表使用者
- [扭曲环面](../twist/overview.md) — ND 策略目标的扭曲圆环几何体
- [返回索引](../index.md)

AllToAll 表 ​

摘要 ​

这在哪里 ​

算法 ​

3. 注册和回读 — 类型 8 / 9 / 0xa ​

用途 ​

算法 ​

功能图 ​

4. MeshNDInfo几何表索引 ​

5. 路由缓冲区 Allocator 值结构 ​