物理核心放置

二进制： extracted/libtpu-0.0.40-cp314-cp314-manylinux_2_31_x86_64/libtpu/libtpu.so(构建 ID 89edbbe81c5b328a958fe628a9f2207d，构建 libtpu_lts_20260413_b_RC00；.text VMA == 文件偏移量 0xe63c000、.rodata VMA == 文件偏移量0x84a0000)。 状态： 重新实现级别 · 证据级别： 已确认(字节锚定)— physical_core_indices 写入路径 (AddCollectivePhysicalCoreIndicesHelper)、读回 (GetPhysicalCoreIndices) 和 tensor_split_mode==2 分类器 (TensorSplitPerCoreClassifier)对所有三个函数的 IDA 反编译进行了交叉检查；下面标记为 [LOW] 的两个剩余子布局 · 第 XIII 部分 — On-Pod Collectives & Barriers / SparseCore-offload Collectives · 返回索引

摘要

本页记录了 SparseCore 卸载的嵌入集合体的 物理核心放置数据路径：集合体的逻辑颜色如何落在一组具体的物理 SparseCore (SC) 核心上，以及如何对每核心 张量分割 (tensor_split_mode == 2) 发射进行键控。它拥有三种字节精确机制：

physical_core_indices 填充 — AddCollectivePhysicalCoreIndices (@0x1c868500)，一个递归异步/融合步行器，及其编写器 AddCollectivePhysicalCoreIndicesHelper (@0x1c868920)，它逐字复制所选核心 absl::Span<long const> (截断的 long→int32)进入每个集合的 *OffloadConfig 变体的 repeated int32 physical_core_indices(原始字段 4)；
physical_core_indices 索引数组布局 + 读回 — [variant+0x18..+0x20] 处的 proto2 RepeatedField<int32>，由 GetPhysicalCoreIndices (@0x1c8692e0) 读回(扩大到 long)到StatusOr<absl::InlinedVector<long,4>>；
tensor_split_mode==2 每核心发射密钥 — TensorSplitPerCoreClassifier (@0x13379de0)，它将每颜色 UniDirRingStrategy 映射到分类密钥 — (axis classcode 0/2/4) | (tensor-split bool) 表示环面轴，或原始 D2DUniDirRingStrategy::core_classifier_ 计数(无split_bit OR)用于 kCoresOnChip — 以及 CanBeCombined (@0x13379d60) 谓词，将每种颜色策略合并到一个每 SC 核心部分张量运算中。

核心选择策略 - 哪些核心符合资格以及 SelectCores 构建所选列表的顺序 - 不在这里拥有；它位于 SC 核心选择(卸载) 和 SC内核选择。 ND 平面/张量分裂因子 推导位于 张量分割/ND 平面。此页面在所选核心 Span 存在时选取并跟踪它进入原型，并在为其键入键的分类器处选取 split-mode-2 发射。

二进制文件中观察到的放置数据路径的合同：

physical_core_indices 是由 SparseCoreQueueAssignment::SelectCores 生成的所选物理 SC 核心 ID 的排序集，由调用者 AssignQueueIDsToAsyncStart (@0x10fdf480) 进行数字化 __sort，然后将其作为absl::Span<long const>。
填充纯粹是副本：AddCollectivePhysicalCoreIndicesHelper 将每个 long 截断为 int32 并将其附加到变体的 RepeatedField<int32> - 没有每个元素的转换、缩放或重新索引。
相同的 Span 驱动两个消费者：它被复制到 physical_core_indices 并附加到 MegaChipParallelismConfig 重复长字段 - 放置和兆芯片并行性共享一个选定的核心列表。
步行者在融合体上递归：步行异步包装融合的 called_computations，并且 {6, 9, 12, 86, 93} 中的每个子集体操作码接收相同的Span。
TensorSplitPerCoreClassifier 返回每个 SC 核心的分组/排序判别式，而不是指令；对于 IciDim::kCoresOnChip，判别式是 D2DUniDirRingStrategy::core_classifier_ 计数 - 巨核跨核分割 (LogicalDevicesPerChip(SparseCore) == 2)，即密集 TensorCore 巨核数据分割的 SC 模拟。

概览

方面	值(字节锚定)
Walker(递归)	`AddCollectivePhysicalCoreIndices` `@0x1c868500` → `bool`
写入器	`…Helper` `@0x1c868920`(匿名)→ `bool`
读卡器	`GetPhysicalCoreIndices` `@0x1c8692e0` → `StatusOr<InlinedVector<long,4>>`
制片人(跨度)	`AssignQueueIDsToAsyncStart` `@0x10fdf480`(呼叫`@0x10fdff30`)
兄弟制片人	`OffloadCollective` `@0x10fc75e0` (`SetMegaChipParallelism` + `AddCollective…`)
索引数组	proto2 `RepeatedField<int32>`位于`[variant+0x18]`标志/`[+0x1c]`尺寸/`[+0x20]` `Rep*`(数据位于`Rep+0x8`)
Oneof 变体指针	AR `[CollectiveOffloadConfig+0x18]` · AG `[+0x20]` · RS `[+0x28]`(判别双字 `[+0x10]`)
Walker操作码集	`{6, 9, 12, 86, 93}`(跳转表 `@0xb438098`，掩码 `0x1240` + 显式 `cmp 0x56/0x5d`)
回读错误行	`523` / `527` / `565` (`backend_config_util.cc`)，所有 `MakeErrorImpl<13>` (内部)
Split-2 分类器	`TensorSplitPerCoreClassifier` `@0x13379de0` → 轴 `0/2/4 \| byte[s+0x43]`； kCoresOnChip → 原始 D2D `core_classifier_`(无 OR)
`ici_dim()`	vtable `+0x68` = `byte[s+0x80] \| 0x100`(`AutoOr`-接合`IciDim`)
组合钥匙	`core_count()`(虚表`+0x30` = `qword[s+0x88]`)通过`CanBeCombined` `@0x13379d60`
读卡器消费者	`CheckCoreAssignmentConsistency` `@0x1c869cc0`

1. `physical_core_indices` 字段 — 阵列布局

physical_core_indices 是 原型字段 4，一个 repeated int32，在每个集合 *OffloadConfig 变体(AllGatherOffloadConfig / AllReduceOffloadConfig / ReduceScatterOffloadConfig)内携带，它们是字节相同的(请参阅**SC-卸载配置生成器** §4)。该字段是标准 proto2 RepeatedField<int32> 并在变体中占据四个偏移量：

偏移(相对变体)	字段	含义
`[variant+0x10]`	`int32` 有位	bit0 = `physical_core_indices` 存在
`[variant+0x18]`	`RepeatedField` 标志 \|竞技场标签	低位 =“有堆 `Rep`”
`[variant+0x1c]`	`int32` `current_size`	元素计数
`[variant+0x20]`	`Rep*`	元素数据从 `Rep+0x8` (`int32[]`) 开始

作者和读者都同意这种布局。在…Helper @0x1c868920中，AR填充清除*(_DWORD*)(v9+28)(= [+0x1c]大小)和*(_BYTE*)(v9+16) &= ~1u(= [+0x10]有位)，然后每个元素存储v16[(int)v14 + 2] = v17(= Rep+0x8 + idx*4)，设置*(_DWORD*)(v9+28) = size+1，并且*(_BYTE*)(v9+16) |= 1u。在 GetPhysicalCoreIndices @0x1c8692e0 中，当堆位 *(_BYTE*)(v9+24) & 1 ([+0x18]) 置位时，读取从 *(_QWORD*)(v9+32) ([+0x20]) 中选取 Rep*，从 *(int*)(v9+28) ([+0x1c]) 中计数，并从Rep+0x8。
注意 — [variant+0x28] 是一个单独的标量，向量归零 (*(_DWORD*)(v_variant+0x28)) 并且填充永远不会接触；它不是 physical_core_indices 重复字段的一部分，该字段占用 [+0x18..+0x20]，如上表所示。

1.1 变体选择之一

physical_core_indices 不是直接到达的顶级 *OffloadConfig 字段 — 它是通过 CollectiveOffloadConfig oneof 包装器到达的。作者和读者都通过 [CollectiveOffloadConfig+0x10] 处的 判别双字 选择变体：

判别位	变体	变体 ptr(相对于 `CollectiveOffloadConfig`)
`& 0x1`	`AllReduceOffloadConfig`	`[+0x18]` (`((_QWORD)cfg+3)`)
`& 0x2`	`AllGatherOffloadConfig`	`[+0x20]` (`((_QWORD)cfg+4)`)
`& 0x4`	`ReduceScatterOffloadConfig`	`[+0x28]` (`((_QWORD)cfg+5)`)
`& 0x8`	`RaggedAllToAllOffloadConfig`	`mutable_ragged_all_to_all_offload_config`
`& 0x10`	`AllToAllOffloadConfig`	`mutable_all_to_all_offload_config`

这两个 *_to_all 变体由模板化 lambda $_0<RaggedAllToAllOffloadConfig> / $_0<AllToAllOffloadConfig> 处理，而不是内联 AR/AG/RS 循环，但字段布局(proto2 RepeatedField<int32> 位于相同的相对偏移量)是相同的。

在…Helper中，判别式为v8 = *((_DWORD*)v6 + 4)(= [cfg+0x10])，测试过& 2(AG，ptr *((_QWORD*)v6+4))，& 1(AR，*((_QWORD*)v6+3))， & 4(RS、*((_QWORD*)v6+5))、& 8(参差不齐→mutable_ragged_all_to_all_offload_config(v6) + $_0)、& 0x10(a2a→mutable_all_to_all_offload_config(v6) + $_0)。 GetPhysicalCoreIndices 选择相同的位(v60 & 2 门控“无集体卸载配置”错误；然后是相同的 & 2 / & 1 / & 4 / & 8 / & 0x10)。

2. 填充路径 — `AddCollectivePhysicalCoreIndices` + `…Helper`

2.1 递归步行器 (`@0x1c868500`)

AddCollectivePhysicalCoreIndices(HloInstruction*, absl::Span<long const>) → bool 不是作者；它是一个递归集体步行器，将相同的 Span 转发到每个子集体：

text

AddCollectivePhysicalCoreIndices(hlo, Span indices):
  CHECK(hlo->IsAsynchronous())                    (@0x1e592520)            // FATAL "inst->IsAsynchronous()" (.cc:433)
  wrapped = hlo->async_wrapped_instruction()       (@0x1e5aa300)
  if opcode_byte[wrapped+0xc] != 0x1b (fusion)                              → return 1
  for each instr in wrapped->called_computations() (@0x1e5885a0):
     switch opcode_byte[instr+0xc]  (jump table @0xb438098, opcode-6 index):
        6 / 9 / 12  → AddCollectivePhysicalCoreIndicesHelper(instr, indices)   // mask 0x1240
        0x56 (86)   → AddCollectivePhysicalCoreIndicesHelper(instr, indices)   // explicit cmp
        0x5d (93)   → AddCollectivePhysicalCoreIndicesHelper(instr, indices)   // explicit cmp
        0x28 (40)   → recurse into instr->called_computations()
        other       → skip
  return 1
```text

操作码集 `{6, 9, 12, 86, 93}` 是 SparseCore 异步集体加上 all-reduce/all-gather/reduce-scatter/all-to-all 系列。操作码 `0x28`(融合)触发递归下降，因此嵌套融合体被完全行走。步行是深度优先的；相同的 `Span` 到达每个匹配的叶子。

> **[已确认 — 符号 + 结构]** 步行器符号 `…AddCollectivePhysicalCoreIndicesEPNS_14HloInstructionEN4absl4SpanIKlEE @0x1c868500` 出现在 `*_functions.json` 中，其两个源调用站点位于 `AssignQueueIDsToAsyncStart` (`@0x10fdff30`) 和 `OffloadCollective` (`@0x10fc871c`)。该条目由 FATAL `CHECK(inst->IsAsynchronous())` (`backend_config_util.cc:433`) 保护 — 非异步输入中止，它**不是**静默无操作。跳转表形式(`@0xb438098`、`0x23` 双字、操作码−6 索引；掩码 `0x1240` = 位 6/9/12；显式 `cmp eax, 0x56` / `0x5d`；`0x28` 的递归目标)从反汇编中进行字节跟踪。 Helper 调用将*相同的* `(rsi=data, rdx=count)` 对转发到每个匹配的指令。

### 2.2 编写者 (`…Helper @0x1c868920`)

```text
AddCollectivePhysicalCoreIndicesHelper(hlo, Span data, count) → bool:
  bc = GetBackendConfig(hlo)            (@0x1c8664a0)         // absent → return 1
  cfg = bc.CollectiveOffloadConfig      ([BackendConfig+0x178], default-constructed if missing)
  select variant by discriminant [cfg+0x10]   (AR/AG/RS/ragged/a2a, §1.1)
  // FILL (byte-identical for AR/AG/RS):
  variant.physical_core_indices.current_size = 0      ([variant+0x1c] = 0)
  variant.has-bit &= ~1                                ([variant+0x10] &= 0xfe)
  for i in 0 .. count-1:
     v = (int32) data[i]                               // movsxd/trunc: long → int32
     RepeatedField<int32>::Add(v)                      // GrowNoAnnotate @0xe68d9e0 if full;
                                                        //   store [Rep+0x8 + size*4]
     variant.physical_core_indices.current_size = i+1  ([variant+0x1c] = i+1)
     variant.has-bit |= 1                              ([variant+0x10] |= 1)
  // FINALIZE:
  cloned = CloneBackendConfigProto(bc)                 (@0x1e60dac0)
  BackendConfigWrapper::operator=([hlo+0x68], cloned)  (@0x1e60de40)
  return 1

填充是逐字复制：唯一的转换是原始字段 4 的 int32 元素类型强制的 long→int32 截断。生产者选择的选定核心 ID 保持不变，按照调用者对它们排序的顺序。

反编译中的AR填充(v8 & 1分支)：*(_DWORD*)(v9+28) = 0、*(_BYTE*)(v9+16) &= ~1u，循环读取v17 = *(_DWORD*)(v11 + 8*v14)(将8字节Span元素截断为4字节存储)、GrowNoAnnotate<…>(v13, v9, …) full，存储v16[(int)v14 + 2] = v17，设置*(_DWORD*)(v9+28) = v18(大小+1)和*(_BYTE*)(v9+16) |= 1u。 AG (v8 & 2) 和 RS (v8 & 4) 分支结构相同(仅变体 ptr 不同)。最终确定：CloneBackendConfigProto((xla*)v44, …)则BackendConfigWrapper::operator=((char*)v53 + 104, …) = [hlo+0x68]。成功的VLOG字符串"backend_config_util::AddCollectivePhysicalCoreIndices( inst, sparse_core_ids) is OK"存在于.rodata中。

3. 回读 — `GetPhysicalCoreIndices` (`@0x1c8692e0`)

text

GetPhysicalCoreIndices(const HloInstruction* hlo) → StatusOr<InlinedVector<long,4>>:
  bc = GetBackendConfig(hlo)
  if !bc                          → MakeErrorImpl<13>("No backend config found", line 523)
  if !(disc & 2 present)          → MakeErrorImpl<13>("No collective offload config found", line 527)
  select the SAME variant by the SAME discriminant bits (§1.1)
  if variant absent / field empty → MakeErrorImpl<13>("No physical core indices found", line 565)
  count = [variant+0x1c]
  Rep*  = [variant+0x20]  (or inline [variant+0x18] when heap-bit clear)
  for each int32 at Rep+0x8 + i*4:  widen int32 → long   (vpmovsxdq, 4-wide)
  build InlinedVector<long,4>:  count <= 4 → inline; >= 9 → exact heap; 5..8 → cap-8 heap
  return StatusOr OK { sret+0=1, sret+8=data, sret+0x10=size, sret+0x18=cap }
```text

读取器与写入器完全相反：相同的一个变体选择，相同的 `[+0x1c]`/`[+0x20]` 字段偏移，将每个存储的 `int32` 扩大回 `long`。结果类型是 `absl::InlinedVector<long,4>` — 在内联缓冲区中停留的计数 ≤ 4；较大的列表将进行堆分配 (`_size_returning_new`)。所有三个错误路径均为 `backend_config_util.cc` 中的 `MakeErrorImpl<13>`(gRPC 代码 13 = INTERNAL)。读取器的使用者是 `CheckCoreAssignmentConsistency` (`@0x1c869cc0`)，它重新验证放置的核心是否符合集体的指令。

> 三个错误字符串+源行号从反编译中精确读取字节：`"No backend config found"`(第**523**行)、`"No collective offload config found"`(第**527**行)、`"No physical core indices found"`(第**565**行)，全部通过`absl::status_internal::MakeErrorImpl<13>(…, "platforms/xla/service/jellyfish/lowering/backend_config_util.cc")`。 `int32`→`long`加宽为`vpmovsxdq` 4宽展开循环； `cmp v12,5`(计数<5 inline) / `v12 >= 9` (exact) split selects the `InlinedVector<long,4>` storage. The OK sret layout (`[+0]=1` tag, `[+8]` data, `[+0x10]` size, `[+0x18]`帽)是字节确认。

---

## 4. `Span` 生产者 — 所选内核的来源

`Span` 填充副本是由 `SparseCoreQueueAssignment::AssignQueueIDsToAsyncStart` (`@0x10fdf480`) 上游一步构建的。完整的**选择策略**——`GetAllowedCores`、五相`SelectCores`贪婪滤波器、成本平局——由**[SC 核心选择(卸载)](sc-core-selection-offload.md)**和**[SC内核选择](../sparsecore/sc-core-selection.md)**拥有；这里只是移交到此页面的填充：

```text
AssignQueueIDsToAsyncStart(hlo):
  mega = GetMegaChipParallelism(hlo)              (@0x1c867b00)  // StatusOr<InlinedVector<long,4>>
  split_axis0 = mega[0] >> 1                       // megacore: 2 cores per chip
  allowed = GetAllowedCores(hlo)                  (@0x10fda3c0)  // btree_set<long> candidate pool
  chosen  = SelectCores(hlo, allowed, …)          (@0x10fdc4e0)  // unsorted, {phase, cost} order
  __sort(chosen)                                  (@0x10fdfde7)  // ASCENDING numeric core ID
  AddCollectivePhysicalCoreIndices(hlo, Span{chosen.data, chosen.size})   (@0x10fdff30)
  // the SAME chosen array is ALSO appended to MegaChipParallelismConfig (repeated long, loop @0x10fdfe80)

此页面的关键观察结果：SelectCores 按 构建顺序 返回核心({相同 ND 平面、数据深度、分配组、非不同平面、后备} × 升序成本)，但调用者 在数字上以 __sort 列表在填充之前。因此，physical_core_indices 始终以升序物理核心 ID 顺序存储，而不是以选择优先级顺序存储 — 选择顺序在内部使用，在原型中不可见。

兄弟制作者 SparseCoreCollectiveOffload::OffloadCollective (@0x10fc75e0) 在 SetMegaChipParallelism (@0x1c867680、@0x10fc86f1) 和 AddCollectivePhysicalCoreIndices 之间配对相同选择的 Span (@0x10fc871c) — 确认放置列表和大型芯片并行列表是同一个数组。

[已确认 — 符号 + 交叉调用] AssignQueueIDsToAsyncStart、GetMegaChipParallelism、SetMegaChipParallelism 和 OffloadCollective 出现在 *_functions.json 中； AddCollectivePhysicalCoreIndices 调用站点 @0x10fdff30(和 @0x10fc871c)和 __sort @0x10fdfde7 是字节跟踪的。 [低] SelectCores 分配的每个核心的确切数字 ID(逻辑颜色 → 物理核心双射/平局决算算法)是选择策略的关注点，并在核心选择页面上记录(作为 5 相贪婪过滤器，而不是封闭式记分器)——此页面只需要输出是排序的 ID 集。

5. `tensor_split_mode==2` 每核发射

当卸载基板采用分裂张量模式(tensor_split_mode == 2，“采用分裂张量模式”。上游路径门控 - 请参阅 张量分割/ND 平面)时，集体的每颜色环作为 每 SC 核心部分张量操作 发出。决定哪些颜色折叠成一个每核操作的分组由 TensorSplitPerCoreClassifier 键入。

5.1 `TensorSplitPerCoreClassifier` (`@0x13379de0`)

text

TensorSplitPerCoreClassifier(UniDirRingStrategy* s) → long:
  split_bit = (byte[s+0x43] != 0)                  // the tensor-split / per-core bool, bit0
  switch (s->ici_dim() & 0x1ff):                   // ici_dim() = vtable+0x68 = byte[s+0x80] | 0x100
     0x100  IciDim::kX           → return 0 | split_bit
     0x101  IciDim::kY           → return 2 | split_bit
     0x102  IciDim::kZ           → return 4 | split_bit
     0x103  IciDim::kCoresOnChip → d2d = dynamic_cast<D2DUniDirRingStrategy*>(s)
                                    CHECK d2d != nullptr  → FATAL streams "Expected
                                       D2DUniDirRingStrategy if the ici_dim is kCoresOnChip" (.cc:156)
                                    c = d2d->core_classifier_   ([d2d+0x90])
                                    CHECK c >= 0     "core_classifier_ >= 0"   (.h:375)
                                    return c                    // NO `| split_bit` on this path
     default → FATAL CHECK "strategy->ici_dim() == IciDim::kCoresOnChip"   (.cc:153)
```text

分类器返回每个 SC 核心的 **key**。对于环面轴变暗(`kX → 0`、`kY → 2`、`kZ → 4` - 每核部分张量环迭代的轴)，低位与张量分割标志 (`class | split_bit`) 进行或运算。对于 `kCoresOnChip`(超级核心跨核心分割 - split-2 数据路径本身)，分类器直接返回原始 `D2DUniDirRingStrategy::core_classifier_` 计数，无需 `split_bit` OR**(反编译中的 `return v9;`)。 `kCoresOnChip` D2D 策略是密集 TensorCore 巨核数据分割的 SC 模拟。

> `@0x13379de0` 的反编译字节精确匹配：`v3 = *((_BYTE*)this + 67) != 0` (`[+0x43]`)；通过`*(_QWORD*)this + 104LL`(vtable `+0x68`)屏蔽的`& 0x1FF`虚拟`ici_dim()`；比较`0x100 → v4=0`、`0x101 → v4=2`、`0x102 → v4=4`，分别达到`LABEL_5: LOBYTE(v2)=v3; return v4 | v2`(`class | split_bit` OR)； `0x103` → `_dynamic_cast(this, typeinfo UniDirRingStrategy, typeinfo D2DUniDirRingStrategy, 0)`，然后 `v9 = v11[18]`(= `[+0x90]`，18×8)与 `CHECK v9 >= 0`，**直接**返回 `v9`(kCoresOnChip 路径在 `split_bit` 中*不*或)。三个诊断字符串 + 源代码行是字节精确的：`"strategy->ici_dim() == IciDim::kCoresOnChip"` (`offload_collective_strategies.cc:153`)； `.cc:156` 流 `"Expected D2DUniDirRingStrategy if the ici_dim is kCoresOnChip"` 处的 D2D 空检查为 `CHECK(d2d_strategy != nullptr)`； `"core_classifier_ >= 0"`(`offload_collective_strategies.h:375`，流媒体 `"Color id is not set."`)。

### 5.2 `UniDirRingStrategy` 现场图

分类器读取的策略字段(来自 `ImplicitUniDirRingStrategy` 构造函数 `@0x1339bca0` 基础和 `D2DUniDirRingStrategy` 构造函数 `@0x1339ba60` 覆盖)：

| 偏移 | 字段 | 注释 |
|--------|-------|-------|
| `[s+0x42]` | `bool` | ctor 布尔值 #1 |
| `[s+0x43]` | `bool` | **张量分割/每核标志**分类器读取 |
| `[s+0x44]` | `byte` `RingDir` |  |
| `[s+0x80]` | `byte` `IciDim` | `0 kX / 1 kY / 2 kZ / 3 kCoresOnChip`； D2D ctor 硬线 `3` |
| `[s+0x88]` | `long` `core_count` | **组合键** (`CanBeCombined`)； D2D 硬连线 `2` |
| `[s+0x90]` | 底座 `bool` · **D2D `long core_classifier_`** | 每核分割计数； D2D ctor 检查 `target.LogicalDevicesPerChip() == 2` (`.h:361`) |
| `[s+0x91]` | `bool` |  |

D2D vtable(`@0x21908db0`，`+0x10` 处的 vptr)解析插槽 `+0x30` → `core_count`(`@0x13399000` = `qword[s+0x88]`)和插槽 `+0x68` → `ici_dim`(`@0x13399020` = `byte[s+0x80] | 0x100`)。 `| 0x100` 是 `AutoOr<IciDim>` 接合位(与 `AutoOr` 打包卸载配置构建器用于 `HierarchicalKind` 的 `AutoOr` 相同 — 请参阅 **[SC-卸载配置生成器](sc-offload-config-builder.md)** §3)。

> **[已确认 — vtable + ctor]** `D2DUniDirRingStrategy` 存在于 `*_functions.json` 中； vtable 插槽重定位(`+0x30 → 0x13399000`、`+0x68 → 0x13399020`)和 D2D ctor 硬连线(`[+0x80]=3`、`[+0x88]=2`、`[+0x90]=` 尾随 `long` arg 和 FATAL `CHECK(target.LogicalDevicesPerChip() == 2)` `offload_collective_strategies.h:361`) 是字节跟踪的。 **[低]** *哪个*调用方计算 `core_classifier_` `long`(即它是否等于 `tensor_split_factor`、`NumScOffloadDevices/LDPC` 或 megacore 2 核计数)未跟踪到构造 D2D 策略的 `AllReduceUnidirNdStrategy::TryCreate` lambda - 分类器直接返回存储的值；它的起源是张量分裂因子问题(**[张量分割/ND 平面](tensor-split-ndplane.md)**)。

### 5.3 联合收割机消费者 — `CanBeCombined` (`@0x13379d60`)

```text
CanBeCombined(a, b) → bool:
  return a->core_count() == b->core_count()                          // vtable +0x30 = [s+0x88]
      && !FLAGS_xla_tpu_impure_coff_never_combine_colors_test_only

§5.1 中的分类器键加上 CanBeCombined 是 DefaultStrategyCombiner<*> 系列(例如 @0x133a4660 / @0x133a3c60)和 MoreRelaxedStrategyCombiner 用于将每种颜色的 UniDirRingStrategy 对象合并到每个 SC 核心的一个部分张量发射：具有相同分类器键和等于 core_count (并且从不组合标志关闭)的颜色折叠在一起。更简单的同级 DimPerCoreClassifier (@0x13379dc0) 仅在张量分割布尔 (byte[s+0x43] != 0) 上键。

CanBeCombined @0x13379d60：v3 = (*(this->vtable+48))(this, a2, a3) 与 (*(a2->vtable+48))(a2)(core_count 虚拟在 +0x30)的比较，由 !FLAGS_xla_tpu_impure_coff_never_combine_colors_test_only 门控(flag-impl 缓存在 qword_2231E730，通过读取) FlagImpl::ReadOneBool)。 DimPerCoreClassifier @0x13379dc0 字面意思是 return *((_BYTE*)this + 67) != 0;(= byte[+0x43])。两者都是字节精确的。 SinglePhaseRSTransferStrategy / SinglePhaseAGTransferStrategy 的 CanBeCombined 重载作为单独的符号出现。 [低] DefaultStrategyCombiner<*> 合并循环主体(密钥实际上如何对颜色进行分组以及将 IciStrategyRingConfig 与合并的每 SC 核心操作发射相结合)由 CanBeCombined/分类器的调用边缘确认，但未扩展到每场发射。

6. 放置 + split-mode-2 — 关系表

数量	源码	角色
选择的物理核心(已排序)	`SelectCores` `@0x10fdc4e0` → `__sort`(呼叫者)	将数组写入`physical_core_indices`
`physical_core_indices` 写入	`AddCollectivePhysicalCoreIndices` `@0x1c868500` + `…Helper` `@0x1c868920`	逐字 `Span`→`RepeatedField<int32>` 副本
索引数组布局	`[variant+0x1c]`=尺寸·`[+0x20]`=`Rep*`(数据`Rep+0x8`)	proto2 `RepeatedField<int32>`(字段 4)
回读	`GetPhysicalCoreIndices` `@0x1c8692e0`	`StatusOr<InlinedVector<long,4>>` (`int32`→`long`)
回读消费者	`CheckCoreAssignmentConsistency` `@0x1c869cc0`	交叉指令核心协议
共享列表	`MegaChipParallelismConfig` 重复长 (`@0x10fdfe80`)	相同选择的阵列 → 兆芯片并行度
每核心类密钥	`TensorSplitPerCoreClassifier` `@0x13379de0`	轴：`(axis 0/2/4) \| split bool`； kCoresOnChip：原始 D2D `core_classifier_`
`ici_dim`	`byte[s+0x80] \| 0x100`(虚表`+0x68`)	`kX` / `kY` / `kZ` / `kCoresOnChip`
D2D 每核分割	`D2DUniDirRingStrategy::core_classifier_` `[s+0x90]` (LDPC(SC)==2)	兆核跨核分割计数
组合钥匙	`core_count()` `[s+0x88]`(虚表`+0x30`)	`CanBeCombined` 分组

7. 验证说明

针对 libtpu.so v0.0.40 (build-id 89edbbe8…) 的 IDA 反编译进行交叉检查：
写入器 (…Helper @0x1c868920) — GetBackendConfig → v43 == 1 门； CollectiveOffloadConfig 判别式 v8 = *((_DWORD*)v6+4) ([cfg+0x10]) 中的一个，具有位 2/1/4/8/0x10 和变体指针 *((_QWORD*)v6+4)/+3/+5 (= [+0x20]/[+0x18]/[+0x28])；清除([+0x1c]=0，[+0x10]&=~1)+逐字long→int32复制循环(GrowNoAnnotate @0xe68d9e0，存储[Rep+0x8+i*4]，[+0x1c]++，[+0x10]|=1)； CloneBackendConfigProto + BackendConfigWrapper::operator=([hlo+0x68]) 回写 — 全部字节精确。
阅读器 (GetPhysicalCoreIndices @0x1c8692e0) — 相同的位； Rep*=[+0x20]，计数=[+0x1c]，数据Rep+0x8； vpmovsxdq int32→long; cmp 5 内联 / >= 9 精确堆 InlinedVector<long,4> 拆分； 523/527/565 行的三个 MakeErrorImpl<13> 字符串； OK sret 布局 +0/+8/+0x10/+0x18 — 所有字节精确。
场偏移协议 - 作者[+0x1c]/[+0x10]和读者[+0x1c]/[+0x18]/[+0x20]独立确认RepeatedField<int32>在[variant+0x18..+0x20]。
分类器 (TensorSplitPerCoreClassifier @0x13379de0) — byte[+0x43] 读取； ici_dim() vtable +0x68 屏蔽 & 0x1FF； 0x100/0x101/0x102 → 0/2/4; 0x103 → dynamic_cast<D2DUniDirRingStrategy> (CHECK != nullptr、.cc:156) + [+0x90] core_classifier_ (CHECK >= 0);默认致命。轴路径返回axis | split_bit； kCoresOnChip 路径返回原始 core_classifier_ (return v9;)，no split_bit OR。源行 153/156 (.cc) 和 375 (.h) 字节精确。
组合 (CanBeCombined @0x13379d60、DimPerCoreClassifier @0x13379dc0) — 通过虚表 +0x30 实现 core_count() 相等，由 FLAGS_xla_tpu_impure_coff_never_combine_colors_test_only 门控； DimPerCoreClassifier = byte[+0x43] != 0 — 都是字节精确的。
符号 — AddCollectivePhysicalCoreIndices、…Helper(+ $_0<Ragged…> / $_0<AllToAll…> lambda)、GetPhysicalCoreIndices、AssignQueueIDsToAsyncStart、OffloadCollective、 GetMegaChipParallelism、SetMegaChipParallelism、D2DUniDirRingStrategy、CheckCoreAssignmentConsistency 和 CanBeCombined 重载均存在于 *_functions.json 中。
[LOW] 由结构/调用边缘确认，但未完全数字解码：
SelectCores 每个候选布局算法(逻辑颜色 → 物理核心双射/平局决胜)——由核心选择页面拥有；这里只有排序的输出重要。
D2DUniDirRingStrategy::core_classifier_ 构造站点(调用者计算存储的 long，以及它是否等于 tensor_split_factor / NumScOffloadDevices 派生 / megacore 2 核计数) - 分类器直接返回它。
DefaultStrategyCombiner<*> 合并循环主体(分类器键如何将颜色分组到发出的每 SC 核心 IciStrategyRingConfig 中)。

交叉引用

SparseCore 卸载放置和选择

SC 核心选择(卸载) — GetAllowedCores 候选掩码 + 提供 SelectCores 的成本/资源模型
SC内核选择 — 生成本页复制的所选核心列表的五相 SelectCores 贪婪滤波器
张量分割/ND 平面 — tensor_split_factor / NumScOffloadDevices + NDPlaneInfo，选择 tensor_split_mode==2 的门

配置 + 基板

SC-卸载配置生成器 — 承载 physical_core_indices 的 *OffloadConfig 结构系列，以及封装 IciDim 读取镜像的 AutoOr
On-Pod 集体 — 剖面图 — 基板分离和 SC 卸载门

同级子系统

分层种类 — AutoOr<bool> 扁平与分层拆分卸载构建器调度
返回索引

物理核心放置 ​

摘要 ​

概览 ​

1. physical_core_indices 字段 — 阵列布局 ​

1.1 变体选择之一 ​

2. 填充路径 — AddCollectivePhysicalCoreIndices + …Helper ​

2.1 递归步行器 (@0x1c868500) ​

3. 回读 — GetPhysicalCoreIndices (@0x1c8692e0) ​

5. tensor_split_mode==2 每核发射 ​

5.1 TensorSplitPerCoreClassifier (@0x13379de0) ​

6. 放置 + split-mode-2 — 关系表 ​

7. 验证说明 ​

交叉引用 ​

SparseCore 卸载放置和选择 ​

配置 + 基板 ​

同级子系统 ​

物理核心放置

摘要

概览

1. `physical_core_indices` 字段 — 阵列布局

1.1 变体选择之一

2. 填充路径 — `AddCollectivePhysicalCoreIndices` + `…Helper`

2.1 递归步行器 (`@0x1c868500`)

3. 回读 — `GetPhysicalCoreIndices` (`@0x1c8692e0`)

5. `tensor_split_mode==2` 每核发射

5.1 `TensorSplitPerCoreClassifier` (`@0x13379de0`)

6. 放置 + split-mode-2 — 关系表

7. 验证说明

交叉引用

SparseCore 卸载放置和选择

配置 + 基板

同级子系统