SPMD 链路计数除数和集体成本模型

地址、vtable 偏移量和 .rodata 常量适用于 libtpu-0.0.40-cp314 轮(构建 ID 89edbbe81c5b328a958fe628a9f2207d)中的 libtpu.so。其他版本有所不同。所有地址均为VMA；在此图像中 .text VMA == 文件偏移量。

摘要

TPU 集体成本模型存在于两个不同的估计器中，它们共享一组 ICI 带宽常数，但回答不同的问题。 SPMD 分区器 估计器 (GetCommunicationTimeInMilliSec) 生成一个挂钟毫秒数字，用于比较候选分片 - 这是 策略/分片选择器 与替代方案进行比较的成本。 捆绑估计器 (CostModel::GetCollectiveCycles) 生成一个 TensorCore 周期数字，存入 23 槽 ResourceVector，调度程序的 MaxResourceCycles 缩减会消耗该数字 - 它是从 GetHloResourcesImpl 获得的五个定价臂之一。

此页面恢复了两个估计器所依赖的三个部分：

1. SPMD ms 公式中的每链路计数除数。它“不是”芯片参数字段，也“不是”平面副本组基数。它是拓扑派生的：虚拟 TpuSpmdPartitioningVisitor::GetCommunicationMultiplier(访问者虚拟表槽 +0x488)返回 ReplicaGroupsOnNDPlane(plane=2).num_mesh_dims + 1 — 集体副本组跨度的 ND 平面环面网格维度数加一。
2. SPMD 挂钟公式本身，time_ms = (bytes/1e9) / (link_count · ici_GBps) · 1000，包括多片 6.0 GB/s 片间覆盖。
3. 完整的每种 GetCollectiveCycles 公式集 — 全聚集、全归约、归约分散、全对全、不规则全对全、集体置换 — 具有每种类型的带宽除数、共享有效带宽项 eff_Bps = ici_GBps · 0.5 · 1e9、seconds→cycles 比例和每环面尺寸 ResourceVector 槽存款 (R[13..18])。

重新实现，合约为：

• 除数标识。 GetCommunicationMultiplier(vtable +0x488，@0x127a16c0)→ ReplicaGroupsOnNDPlane(…, plane=2, false).num_mesh_dims + 1，默认为1，无DeviceAssignment，多片单组短路。
• SPMD ms 公式 以及 6.0 GB/s DCN/OCS 覆盖和 1e9 / 1000.0 常数。
• 共享的 eff_Bps 减半(×0.5 = 双向 ICI 环的每个方向)和 freq_MHz · 1e6 循环比例(编码为 · 1000.0 两次)。
• HloOpcode 调度(跳转表 @0xae0e9e0，索引 = opcode − 6)以及异步 shell 和广播的零成本默认值。
• 每个种类带宽除数(全聚 1D ÷2 / 2D ÷4；全归约 ÷2·num_dims；全对全 ÷links/per_link；集体置换 ÷1)和 R[13..18] ICI 插槽存款规则。

SPMD除数(vtable +0x488)	TpuSpmdPartitioningVisitor::GetCommunicationMultiplier @ 0x127a16c0
除数身份	ReplicaGroupsOnNDPlane(plane=2).num_mesh_dims + 1(拓扑衍生)
SPMD ms 估计器	TpuSpmdPartitioningVisitor::GetCommunicationTimeInMilliSec @ 0x127a19e0
束周期估计器	CostModel::GetCollectiveCycles @ 0x130abfc0
操作码调度	跳转表@0xae0e9e0，索引=opcode − 6
有效带宽	eff_Bps = IciGigabytesPerSecond() · 0.5 · 1e9
ICI ResourceVector 插槽	R[13..18] = Ici{Y,X,Z}{Plus,Minus}(3 个环面变暗 × 2 个方向)
循环转换	cycles = volume/divisor · TC_freq_MHz · 1e6
常数	0.5 @ 0xa2df5c8 · 1e9 @ 0xa2de620 · 1000.0 @ 0xa2e0430 · 6.0 @ 0xa2de720 · 4.0 @ 0xa2de830

---

第 1 部分 — SPMD 链路计数除数

槽位分辨率

SPMD 分片估计器 TpuSpmdPartitioningVisitor::GetCommunicationTimeInMilliSec (@0x127a19e0) 将字节量除以通过访问者 this 上的虚拟调用读取的两项的乘积：

• IciGigabytesPerSecond() — 每代聚合 ICI 带宽(Target vtable +0x5d8，反编译为 (*(…)(*this + 1496))(this))。
• 链接计数除数 — 访问者自己的 vtable 槽 +0x488。

+0x488 的 Itanium-ABI 分辨率：vtable for TpuSpmdPartitioningVisitor 符号位于 0x218e2508；该对象的vptr是sym + 0x10 = 0x218e2518；插槽+0x488位于0x218e29a0；存储在那里的 R_X86_64_RELATIVE reloc 指向 0x127a16c0 = TpuSpmdPartitioningVisitor::GetCommunicationMultiplier(CollectiveDeviceListBase const&)。相邻插槽确认布局：+0x480 → 0x127a19e0(ms 估计器本身)和 +0x478 → 0x127a15e0 (GetComputationTimeInMilliSec)。

注意 — 除数是访问者上的虚拟方法，而不是 Target 上的虚拟方法。此版本中没有 Target::NumIciLinks 访问器；每链路带宽(ICIPerLinkDataRate，每方向~2-SerDes 数字)是一个单独的带宽常量，而不是这个计数。两者在公式中相乘，但在不同的地方计算。请参阅 ICI 概述。

GetCommunicationMultiplier 计算什么

@0x127a16c0 的反编译很短且字节精确：

uint GetCommunicationMultiplier(visitor *this, CollectiveDeviceListBase &devList) {
    if (!this->device_assignment /* this+0x568 */)              // no DeviceAssignment
        return 1;                                               // default multiplier
    if (Target::GetMultiSliceTopology(this->target /* +0x558 */)) {
        // MULTI-SLICE: build the cross-slice transfer groups (mode 3)
        g = TransferStrategy::ConstructSliceTransferGroup(mode=3, devList, …);
        if (single_cross_slice_group)                           // g has one group
            return 1;                                           // one group ⇒ 1×
        localize → ToReplicaGroup(mode=3) →
        nd = ReplicaGroupsOnNDPlane(target, devAssign, devList, plane=2, false);
        return (uint8)nd.result_byte + 1;                       // [A]
    } else {
        // SINGLE-SLICE:
        nd = ReplicaGroupsOnNDPlane(target, devAssign, devList, plane=2, false);
        return (uint8)nd.result_byte + 1;                       // [B]
    }
}
```text

在 IDA 列表中为：`v2 = 1; if (!*((_QWORD*)this + 173)) return v2;`(字段 `+0x568` = `173·8`)；单片臂称为`ReplicaGroupsOnNDPlane(&v26, target, devAssign, a2, 2, 0)`，然后是`v7 = v27[8]; return v7 + 1`；多切片臂调用`ConstructSliceTransferGroup(…, 3, …)`，当组计数检查失败时返回`1`，否则通过`ReplicaGroupsOnNDPlane(…, 2, 0)`→`v7 = v24[16]; return v7 + 1`重新推导。两个分支都传递 `plane = 2` 和 `bool = false`，获取结果结构的第一个字节 (`movzbl`)，并加一。

> **除数标识：** `link_count = ReplicaGroupsOnNDPlane(plane=2).num_mesh_dims + 1`。

### 为什么是拓扑衍生的

`ReplicaGroupsOnNDPlane` (`@0x1c890960`) 返回一个小结构 `{ok_byte, vector<MeshNDInfo>}`，在 `NDPlaneCacheKey` 上键入的 `NodeHashMap` 中存储。它将集体的副本组分解为物理环面：

- 将每个逻辑设备 ID 映射到其物理 TensorCore 位置(`TensorCoreLocationForLogicalDeviceId` @ `0x1c8904e0`)，
- 读取每个设备的物理芯片坐标(`TpuCoreLocation::Chip`@`0x20ad6720`)，
- 为组跨度的每个环面网格*维度*构建一个 `MeshNDInfo`(`ReplicaGroupsOnNDPlaneImpl::$_0` @ `0x1c896400`；该实现将生成的向量大小与 `2` 进行比较)。

返回的 `result_byte` 是**网格尺寸计数** - 集体的组延伸穿过多少个环面轴。 `+1` 在 `1×` 带宽上铺设一个 0 维/简并群。

语义解释(“副本组跨度的 ND 平面网格尺寸数”)是从工作器的结构中读取的，而不是从标签中读取的； *byte = mesh-dim 计数，则 `+1`* 是字节精确的，并且芯片坐标调用链被完全跟踪。 `MeshNDInfo` 的精确逐场布局和内部维度折叠算法未解码(参见 [置信度和未清项目](#confidence--open-items))。

### SPMD挂钟公式

`GetCommunicationTimeInMilliSec` (`@0x127a19e0`) 组装：

```c
ici_GBps   = target.IciGigabytesPerSecond();                 // Target vtable +0x5d8
if (multi_slice && single_cross_slice_transfer_group)
    ici_GBps = 6.0;                                          // const 0xa2de720 — DCN/OCS inter-slice rate
link_count = this->GetCommunicationMultiplier(devList);      // visitor vtable +0x488
           = ReplicaGroupsOnNDPlane(plane=2).num_mesh_dims + 1;
time_ms    = (bytes / 1e9) / (link_count * ici_GBps) * 1000.0;
//           const 0xa2de620 = 1e9 ; const 0xa2e0430 = 1000.0

反编译尾部为 vdivsd ÷0xa2de620(=1e9)、vmulsd ×var_30(=link_count·ici)、vdivsd、vmulsd ×0xa2e0430(=1000.0) — 字节→GB，除以乘积，毫秒。 0xa2de720 = 6.0多切片覆盖仅出现在该估计器中；捆绑估计器从不读取它。

明白 — 仅当多切片拓扑上存在单个跨切片传输组时，6.0 覆盖才会触发 — 即 DCN/OCS 切片间链路速率，远低于切片内 ICI 速率。当集合跨越多个跨切片组时，应用常规 ReplicaGroupsOnNDPlane(plane=2)+1 除数。在多切片上始终使用 6.0 的重新实现会对多组跨切片集合进行错误定价。

功能	地址	角色
GetCommunicationTimeInMilliSec	0x127a19e0	SPMD 分片 ms 估计器；公式 + 6.0 覆盖
GetCommunicationMultiplier	0x127a16c0	vtable +0x488 除数 = ReplicaGroupsOnNDPlane(plane=2)+1
ReplicaGroupsOnNDPlane	0x1c890960	plane=2 网格数；已记忆 NDPlaneCacheKey
ReplicaGroupsOnNDPlaneImpl::$_0	0x1c896400	每个副本组 mesh-dim 工作线程(vec 大小与 2)
TensorCoreLocationForLogicalDeviceId	0x1c8904e0	逻辑 id → 物理 TC 位置
TpuCoreLocation::Chip	0x20ad6720	物理芯片坐标
ConstructSliceTransferGroup	0x14b8ca20	多片跨片传输组(模式3)
Target::GetMultiSliceTopology	0x1d617980	多片拓扑探头

---

第 2 部分 — 捆绑成本模型：GetCollectiveCycles

CostModel::GetCollectiveCycles (@0x130abfc0) 是 bundle 集体估计器。它是从 GetHloResourcesImpl 的集​​体臂到达的(IsSupportedCollectiveHlo → GetCollectiveCycles，作为标量循环加上每链路 ResourceVector 插槽存入 rv[+8])。与 SPMD 估计器不同，它对**片内每环面维度 ICI 环进行建模，并且从不应用 6.0 多片覆盖。

共享有效带宽项

该功能通过读取每代ICI率并将其减半来打开：

ici_GBps = target.IciGigabytesPerSecond();   // Target vtable +0x5d8 (1496) @0x130abfe3
eff_Bps  = ici_GBps * 0.5 * 1e9;             // const 0xa2df5c8 = 0.5 ; 0xa2de620 = 1e9
// seconds → cycles scale = TC_freq_MHz * 1e6, encoded as (· 1000.0) twice (0xa2e0430)
```text

反编译准确地显示了这一点：紧随 vtable-`+0x5d8` 调用之后的 `vmulsd xmm0, qword_A2DF5C8` (×0.5)，然后是 `vmulsd xmm1, qword_A2DE620` (×1e9)。每个种类分支将秒转换为周期，并应用两次 `TensorCoreFrequencyInMegaHertz()` 和 `vmulsd …, 0xa2e0430` (`·1000·1000 = ·1e6`)。

> **注意 —** **×0.5** 减半将 ICI 环建模为双向，仅对每个方向的聚合带宽进行计费。在任何集合分支中**没有附加延迟项** - 没有 `InitialDmaLatency` 调用，没有固定 ns 添加。捆绑包的集体成本是纯带宽，以 TensorCore 周期表示。

### 操作码调度

在 `eff_Bps` 设置之后，调度是一个 `switch(opcode)` 编译为 `.rodata 0xae0e9e0` 处的自相关跳转表，索引为 `opcode − 6`(反编译中为 `v17 = (uint)(v15 - 6)`)。 `HloOpcode` 整数通过 `HloOpcodeString` 长度表进行字节精确验证。

| 操作码 | 名称 | 分公司 |
|---:|---|---|
| 6 | `all-gather` | AllGather 分支 @ `0x130ac06c` |
| 8 | `all-gather-start` | AllGather 分支 @ `0x130ac06c` |
| 9 | `all-reduce` | AllReduce分支@`0x130ac14c` |
| 11 | `all-reduce-start` | AllReduce分支@`0x130ac14c` |
| 12 | `all-to-all` | 尾调用 `ComputeAllToAllCycles` @ `0x130ae8e0` |
| 34 | `collective-permute` | CollectivePermute 分支 @ `0x130ac40f` |
| 36 | `collective-permute-start` | CollectivePermute 分支 @ `0x130ac40f` |
| 86 | `ragged-all-to-all` | 尾调用 `ComputeRaggedAllToAllCycles` @ `0x130aea80` |
| 93 | `reduce-scatter` | AllReduce-族路径 (`ComputeAllReduceCycles`) |
| 7, 10, 33, 35 | `-done` / `collective-broadcast` / `cp-done` | 默认 @ `0x130ae546` ⇒ **0 周期** |

> **GOTCHA —** 异步 shell (`-start` / `-done`) 和 `collective-broadcast` 贡献零 ICI 捆绑成本；成本根据*数据携带*操作码收取。操作码 7/10/13..33/35 所有 `goto` 都是零成本标签。对 `-start` 和 `-done` shell 双重计数进行收费的重新实现。

### 每环面维度 ICI 插槽

`GetResourceFromIciResource` (`@0x1c894c00`) 将 `IciResource` 枚举映射到 `ResourceVector` 槽索引。反编译返回有效资源的 `v1 + 0x10000000D` — 即低位双字中的基槽 `0xD` (= 13) 加上高位双字中的 `0x100000000`“存在”标志。调用者测试 `& 0x100000000` 的有效性并使用低位 dword 作为插槽：

| 插槽 | 偏移 | 名称 | 圆环轴/方向 |
|---:|---|---|---|
| `R[13]` (`0xd`) | `+0x68` | `IciYPlus` | 暗淡 Y, + |
| `R[14]` (`0xe`) | `+0x70` | `IciYMinus` | 暗淡 Y, − |
| `R[15]` (`0xf`) | `+0x78` | `IciXPlus` | 暗 X，+ |
| `R[16]` (`0x10`) | `+0x80` | `IciXMinus` | 暗淡 X, − |
| `R[17]` (`0x11`) | `+0x88` | `IciZPlus` | 暗 Z，+ |
| `R[18]` (`0x12`) | `+0x90` | `IciZMinus` | 暗淡 Z，- |

这些正是 [资源枚举(23 槽)](../cost/resource-enum.md) 中命名的六个 ICI 环链路插槽：三个环面尺寸 × 两个环方向。每个种类的分支都通过 `ResourceVector::Acc(a3, slot, cycles)` 循环存入这些插槽。当环面轴降级时，[退化轴路径](degraded-axis.md)(`GetDegradedAxis` @ `0x1c894c20`、`Target::Is{X,Y,Z}Degraded` @ `0x1d615940`/`0x1d615960`/…)会删除尺寸的两个槽。

---

## 第 3 部分 — 每类公式

表示法：`B` = 通过 `GetShapeSize` (`@0x130aec20`) 的字节量； `eff_Bps` 如第 2 部分； `cyc(x) = x · TC_freq_MHz · 1e6`。存款落在针对集体拓扑维度有效的 `R[13..18]` 插槽中。

### 全聚集(操作码 6 / 8) — `@0x130ac06c`

```c
n = LogicalDeviceCount ratio = out_shape_size / in_shape_size;   // idiv @0x130ac141
B = (n - 1) * output_shape_size;                                 // (n-1) shards transferred
if (UseAllGather2D)   c = cyc(B / (4 * eff_Bps));                // const 0xa2de830 = 4.0
else                  c = cyc(B / (2 * eff_Bps));                // 1D bidirectional ring (×2)
// deposit c into R[13/14], R[15/16], R[17/18] gated by topology dim bitmask &1/&2/&4

AllGatherEmitter::UseAllGather2D(@0x13801740)选择路径：2D环使用2个环面dim × 2个方向= 4个并发链路(除数4)，1D环使用1个dim的2个方向(除数2)。反编译时2D路径乘以qword_A2DE830(4.0)；一维路径为 vaddsd xmm1,xmm1,xmm1 (×2)。存款由 & 1 → Acc(13)+Acc(14)、& 2 → Acc(15)+Acc(16)、& 4 → Acc(17)+Acc(18) 控制。

全归约(操作码 9 / 11) — @0x130ac14c

B = 2 * operand_shape_size;                          // ReduceScatter + AllGather = 2× volume
if (DoesReplicaGroupFormNDPlane(...)) {               // fills NDTopologyInfo dim bitmask
    num_dims = popcnt(NDtopo_dims & 7);               // active torus axes (__popcnt(&7))
    c = cyc(B / (2 * num_dims * eff_Bps));
    // deposit c into the 2 slots of EACH active dim (test &1→13,14 ; &2→15,16 ; &4→17,18)
} else {                                              // cross-module / non-ND-plane fallback
    c = ComputeAllReduceCycles::$_0(...);             // @0x130d0040
    //  B' = AllReduce.size_field ; c = cyc(B' / (2 * eff_Bps))   // single bidirectional ring
    //  links = EstimatePhysicalLinksUsed(...) ; deposit into all 6 slots by resource
}
```text

`IsCrossModuleReduceInstruction`(`@0x137dac40`)和`DoesReplicaGroupFormNDPlane`(`@0x1c88bfc0`)决定路径；反编译显示`v460 = 2 * ShapeSize`、`__popcnt((v465)&7)`，然后是`vmulsd …, var_58(eff_Bps)`、`vdivsd`。跨模块后备委托给 `ComputeAllReduceCycles::$_0` lambda，它使用 `EstimatePhysicalLinksUsed` 并在所有六个插槽中存款。

### 减少分散(操作码 93)— AllReduce 系列

在 AllReduce 路径中处理(`opcode == 93` 形状比臂和 `0x133e1740` 处的 SparseCore `GetCollectiveOffloadConfig` 探针)。它被建模为 all-reduce 分解的 **reduce-scatter 阶段**：具有相同 `eff_Bps / num_dims` 项的相同 per-dim 环沉积 — all-reduce 的一半成本双胞胎(all-gather 阶段是互补的一半)。没有单独的跳转表槽。

### All-to-all(操作码 12)和 Ragged-all-to-all(操作码 86)

均尾调用至 `ComputeAllToAllCyclesHelper` (`@0x130d02a0`)；不规则变体设置了 `bool` 标志，但共享公式：

```c
B        = operand_shape_size * group_size;          // imul
links    = EstimatePhysicalLinksUsed(...);           // @0x1c8939c0 — |sorted distinct IciResource set|
per_link = per_link_table[is2D];                     // table {1D→2.0, 2D→4.0}
c        = cyc(B * per_link / links / eff_Bps);      // vmul [tbl+is2D*8] ; ÷links ; ÷eff_Bps
// deposit c into ALL 6 slots R[13..18] — all-to-all saturates every ICI link

反编译：vmulsd xmm2, qword [rax+rcx*8](is2D索引的每个链接表)，vdivsd xmm1, xmm2, xmm1(÷links)，vdivsd xmm0, xmm0, xmm1(÷eff_Bps)，其中EstimatePhysicalLinksUsed称为上面。 4.0 常量 (0xa2de830) 显示为 2D 表条目。

集体置换(操作码 34 / 36) — @0x130ac40f

B = GetShapeSize(operand);
c = cyc(B / eff_Bps);                                // single-direction point-to-point (NO ×2)
if (AllPairsUseSameIciLink(target, devAssign, source_target_pairs, channel_id)) {
    r = GetResourceFromIciResource(resource_id);     // narrow to ONE ICI slot
    Acc(a3, r, c);
} else {
    Acc(a3, 13, c); Acc(a3, 14, c); Acc(a3, 15, c); Acc(a3, 16, c); Acc(a3, 17, c); Acc(a3, 18, c);
}
```text

集体置换是唯一一种单独除以 `eff_Bps` 的类型(无 `×2` 因子)——沿一个环方向的定向点对点传输。 `AllPairsUseSameIciLink` (`@0x1c88de40`) 检测每一对 `(src,dst)` 是否乘坐一条 ICI 链路；如果是这样，押金将缩小到 `GetResourceFromIciResource` 返回的单个资源槽，否则成本将分摊到较低的 ICI 槽。

---

## 每种配方汇总

| 集体(操作码) | 卷 `B` | 带宽除数 | 目录 | 存入槽位 |
|---|---|---|---:|---|
| 全聚 (6/8) 1D | `(n−1)·out_size` | `2 · eff_Bps` | 2 | 每暗 `R[13..18]` |
| 全聚 (6/8) 2D | `(n−1)·out_size` | `4 · eff_Bps` | 4 | 每暗淡(2 暗淡) |
| 全缩小 (9/11) ND 平面 | `2 · operand_size` | `2·num_dims · eff_Bps` | 2/暗淡 | 每个主动调光(每个 2 个) |
| 全归约跨模块 | `AR.size_field` | `2 · eff_Bps` | 2 | 全部 6 个(由 `links` 提供) |
| 减少分散 (93) | `operand_size`(RS相) | `2·num_dims · eff_Bps` | 2/暗淡 | 每个主动调光 |
| 全对全 (12) | `op_size·group_size` | `links / per_link · eff_Bps` | 变量 | 全部 6 `R[13..18]` |
| 衣衫褴褛 (86) | `op_size·group_size` | `links / per_link · eff_Bps` | 变量 | 全部 6 |
| 集体排列(34/36) | `operand_size` | `1 · eff_Bps` | 1 | `R[13/15/17]±` 或 1 个插槽 |
| `*-done` / 广播 / cp-done | — | — | — | 0 周期(默认) |

`eff_Bps = IciGigabytesPerSecond · 0.5 · 1e9`; `cycles = volume/divisor · TC_freq_MHz · 1e6`。 `num_dims = popcnt(active torus axes)`； `links = EstimatePhysicalLinksUsed`； `per_link = {1D→2.0, 2D→4.0}`。

| 功能 | 地址 | 角色 |
|---|---|---|
| `CostModel::GetCollectiveCycles` | `0x130abfc0` | 束集体估计器； `eff_Bps` + 操作码调度 |
| `ComputeAllToAllCycles` | `0x130ae8e0` | 操作码 12 条目 |
| `ComputeRaggedAllToAllCycles` | `0x130aea80` | 操作码-86 条目 |
| `ComputeAllToAllCyclesHelper` | `0x130d02a0` | `B·per_link/links/eff_Bps`；每链接表 `{2,4}` |
| `ComputeAllReduceCycles::$_0` | `0x130d0040` | 跨模块AR回退；全6槽存款 |
| `GetResourceFromIciResource` | `0x1c894c00` | `IciResource → slot 0xD..0x12` + 当前标志 |
| `EstimatePhysicalLinksUsed` | `0x1c8939c0` | 拓扑链路计数器(芯片坐标行走) |
| `AllGatherEmitter::UseAllGather2D` | `0x13801740` | 1D ÷2 与 2D ÷4 选择 |
| `DoesReplicaGroupFormNDPlane` | `0x1c88bfc0` | 填充 `NDTopologyInfo` 暗淡位掩码 |
| `AllPairsUseSameIciLink` | `0x1c88de40` | 集体置换单链路窄化 |
| `GetShapeSize` | `0x130aec20` | 字节量`B` |

---

## 置信度和未清项目 {#confidence--open-items}

字节精确、反编译确认的核心(本页的主要声明)：

- vtable `+0x488`→`GetCommunicationMultiplier` reloc解析及其控制流程(默认`1`，no-DeviceAssignment短路，多片单组`1×`，`ReplicaGroupsOnNDPlane(plane=2).result_byte + 1`)。
- SPMD ms 公式、`6.0` 多切片覆盖和 `1e9` / `1000.0` 常数。
- `eff_Bps = ici · 0.5 · 1e9` 术语、`opcode − 6` 跳转表、每个种类带宽除数(AG 1D÷2 / 2D÷4、AR ÷`2·num_dims`、全部到全部 ÷`links/per_link`、CP ÷1)以及`R[13..18]` 槽存款。
- `HloOpcode` 整数 → 所有集合操作码的名称映射(长度验证)。

通过公式形状而不是命名常量确认的项目，以及未现场解码的项目：

- 精确的整数 `ReplicaGroupsOnNDPlaneImpl::$_0` 作为“平面计数”返回 — 被证明是 `vector<MeshNDInfo>` 大小(与 `2` 相比)和 `+1`'d，但每个字段的 `MeshNDInfo` 编码(`MeshDim`步幅/大小/原点)和维度折叠算术未逐场解码。 **高** 在计数语义上，字段布局是开放的。
- `EstimatePhysicalLinksUsed` 链接计数：**已解决** — 该函数构建定向 `IciResource` 的*去重、排序集*(`1..6` = 3 个环面变暗 × 2 个方向；`absl` 哈希集插入 + `std::__introsort<IciResource*>`反编译)，除数是集合基数 `|EstimatePhysicalLinksUsed(...)|`，**不是**每维范围的乘积。看到了多片跨片分支(`EstimatePhysicalLinksUsed::$_0` @ `0x1c894ac0`，全局→本地逻辑ID解析器)，但未扩展。 **高**关于集合基数语义；参见 [SC侧扭转§3](../twist/sc-side-twist.md#3-estimatephysicallinksused--the-physical-ici-link-estimator)。
- `{0,1} = {2.0, 4.0}`之外的每链接表：索引≥2解码为垃圾(仅执行1D/2D)； 3D all-to-all 是否使用真正的索引 2 条目尚未测试。 **低** 假设的 3D 条目。
- `AllPairsUseSameIciLink` 资源 ID 源自 `source_target_pairs` 几何(1-slot-vs-3-slot 谓词)。 **高** 在现有分支上，几何映射已打开。
- SparseCore `GetCollectiveOffloadConfig` (`@0x133e1740`) 用于卸载的reduce-scatter/all-reduce 的备用成本重新路由：读取探针，但其备用路径未单独解码。探针上**高**，卸载成本路径打开。请参见 [SC 卸载配置](sc-offload-config-builder.md) / [SC磁芯选择](sc-core-selection-offload.md)。

---

## 交叉引用

- [集体概述](overview.md) - 该成本模型服务的策略选择器和集体算法系列
- [SelectNDStrategy(选择器)](strategy-nd-picker.md) — *消耗*这些成本的分片/算法决策； `GetCommunicationTimeInMilliSec` ms 图是它比较的成本与成本信号
- [降级轴摄取](degraded-axis.md) — `GetDegradedAxis` / `Is{X,Y,Z}Degraded` 从存款中删除环面轴的两个 ICI 槽
- [GetHloResources路由](../cost/gethloresources-routing.md) — 调度到 `GetCollectiveCycles` 的 `GetHloResourcesImpl` 集体臂
- [资源枚举(23 槽)](../cost/resource-enum.md) — `ResourceVector` 和 `R[13..18]` `Ici{Y,X,Z}{Plus,Minus}` 插槽此成本模型存入
- [成本模型概述](../cost/overview.md) — 每代 `Performance` / `CycleTable` 系列以及这些公式使用的 `TensorCoreFrequencyInMegaHertz` 循环转换
- [ICI 概述](../ici/overview.md) — 物理 ICI 环链路，其单向带宽为 `×0.5` 因子模型
- [扭曲概述](../twist/overview.md) — 为 `ReplicaGroupsOnNDPlane` / `NDTopologyInfo` 提供数据的扭曲环面复制对几何结构