SC 核心选择
本页中的每个地址、偏移、字符串和常量都从
libtpu-0.0.40-cp314wheel 中的libtpu.so逐字节读取（BuildID md589edbbe81c5b328a958fe628a9f2207d，buildlibtpu_lts_20260413_b_RC00）。.textVMA == file offset（0xE63C000）；.rodataVMA == file offset（0x84A0000）。其他版本会不同。
摘要
本页描述 SparseCore 后端在把 embedding/collective async op 分配给一组具体物理 SparseCore 核心时运行的物理核心选择策略。它由 xla::jellyfish::SparseCoreQueueAssignment 的两个成员函数拥有：GetAllowedCores（0x10FDA3C0）为一个 collective 计算候选核心 mask，SelectCores（0x10FDC4E0）把该 mask 转换成一个有序核心列表。产物是 absl::Span<long const>，调用方会对其进行数值排序，并通过 AddCollectivePhysicalCoreIndices（0x1C868500）写入 op 的 physical_core_indices backend-config 字段。
重新实现者需要的心智模型是，SelectCores 不是数值 scorer —— 在其 0x28C1 字节主体（nm -S）的任何位置，都没有闭式 score(core)，也没有 core→queue 双射算术。它是一个确定性的贪婪优先级过滤器：它把允许的 btree-set 物化为 flat vector，将该 vector 一次性排序为逐核心 cost 升序（这是 double cost 参数唯一被消费的位置），然后对 cost 排序后的候选运行五个有序 pass。每个 pass 在候选第一次满足该 pass predicate 时追加它，而 predicate 会针对已经分配的 async collective 的运行中 vector<Info> 求值。因此候选的最终 rank 是 (pass index 1..5) 为主、(ascending cost) 为次。五个 predicate 依次为：（P1）核心已经在同一个 ND plane 上运行 collective；（P2）持有该核心的已分配 op 与当前 op 存在数据依赖（通过 HloReachabilityMap 探测）；（P3）该 op 是预定 assignment group 的成员，且该 group 其他成员持有该核心；（P4）核心没有在不同 ND plane 上运行 collective；（P5）fallback —— 追加所有剩余的允许核心。P5 保证所选集合覆盖整个 allowed set；P1–P4 只确定顺序，调用方按 device/megachip 数量截断后，该顺序决定成员。
GetAllowedCores 从两个输入构建候选 mask，并通过两个 Swiss table 求交：megachip 逐轴 chip ID（GetChipIDsFromParallelismConfig → GetMegaChipParallelism）和逐 collective scheduling resource-type ID {0, 23, 24, 25, 26, 27, 28}（GetSparseCoreResources：对 offload-config collective-type enum 的 7 分支 switch，其中六个分支插入固定常量，enum-4 分支则返回 unwrapped async op root opcode 上的 AsyncTracker::GetResourceTypeForOp）。逐 resource 的线程本地 reservation budget（通过 __tls_get_addr 读取，逐次分配递减，由 >= 2 gate）会排除该 resource 的 reservation 已耗尽的核心 —— 这是 embedding-device / reserved-core 效应。输出是允许的 SC core ID 的 btree_set<long>。
本页分为四个单元：端到端流水线（选择位于候选 mask 和 proto 写入之间的位置）、GetAllowedCores（候选 mask 构建和 reservation budget）、SelectCores（五阶段过滤器、Info struct、cost tie-break），以及 tensor_split_factor 交互（所选核心的数量如何被下游 collective ring strategy 消费）。约束逐 chip 核心数量的 4:1 SC:TC 比例见 SparseCore Hardware Architecture。
对重新实现而言，契约是：
- 选择是
physical_core_indices的两阶段 producer，而不是单个 scorer。GetAllowedCores（mask）→SelectCores（order）→ 调用方数值__sort→AddCollectivePhysicalCoreIndices。SelectCores返回其 vector 时未排序；数值排序发生在调用方。 SelectCores是五阶段贪婪优先级过滤器。 Pass 顺序为{same-plane, data-dep, assign-group, not-different-plane, fallback}；在一个 pass 内，cost 升序stable_sort是 tie-break。成员由已经分配的 collective 状态决定，而不是由任何逐候选数字决定。- 唯一的 "cost" 是 tie-break，不是 placement metric。
double cost参数供给一个逐核心权重，stable_sortcomparator（$_0，vucomisd）读取它来给同优先级候选排序。较低 cost 先追加。 GetAllowedCores是 resource-reservation 交集。 允许核心 = 由逐 collective resource-type ID{0,23..28}分组的 megachip chip ID，减去逐 resource 线程本地 reservation budget 耗尽的核心。tensor_split_factor是 consumer，不是选择输入。 它位于 collective offload-config 中，控制 ring strategy 如何分区所选核心；该 strategy 强制color_strategies_size() % tensor_split_factor == 0。它不会在SelectCores或GetAllowedCores内部分支。
| 候选 mask | SparseCoreQueueAssignment::GetAllowedCores(HloInstruction*)（0x10FDA3C0）→ btree_set<long,…,256> |
| 有序选择 | SparseCoreQueueAssignment::SelectCores(hlo, allowed, devcount, cost, assigned, reach, assign_groups)（0x10FDC4E0）→ StatusOr<vector<long>> |
| Mask 输入 | GetChipIDsFromParallelismConfig（0x10FDBF40）∩-by GetSparseCoreResources（0x10FDC0A0） |
| Resource-type 集合 | {0, 23, 24, 25, 26, 27, 28}（.rodata 0xAC0A8E0..0xAC0A910，逐字节读取；jump table @0xAC0A82C） |
| Reservation budget | 逐 resource 线程本地 long（__tls_get_addr(&qword_22048D78)），递减，由 >= 2 gate |
| 选择 pass | P1 same-plane · P2 data-dep（HloReachabilityMap）· P3 assign-group · P4 not-different-plane · P5 fallback |
| Tie-break | cost 升序 stable_sort comparator $_0（vucomisd，0x10FE8AE0） |
| 调用方 / sink | AssignQueueIDsToAsyncStart（0x10FDF480）→ __sort → AddCollectivePhysicalCoreIndices（0x1C868500）→ physical_core_indices |
端到端流水线
目的
核心选择是 queue assignment 的一个子步骤。对每个 async-start collective，SparseCoreQueueAssignment pass 必须决定它占用哪些物理 SC core；该决策成为一个 long core ID 列表（截断为 int32），存储在 op 的 physical_core_indices backend-config 字段中，下游 emitter 会读取它。选择本身从不接触 proto —— 它产生一个 absl::Span<long const>，由调用方提交。
入口点
调用方 AssignQueueIDsToAsyncStart（0x10FDF480）中的链条已逐字节确认：
AssignQueueIDsToAsyncStart (0x10FDF480) ── per-collective driver
├─ GetAllowedCores(hlo) (0x10FDA3C0) ── candidate btree_set<long> [call @ caller line ~236]
├─ SelectCores(hlo, allowed, devcount, cost, ── ordered vector<long> (UNSORTED) [call @ caller line ~406]
│ assigned, reach, assign_groups) (0x10FDC4E0)
├─ std::__u::__sort<__less<long>>(begin, end) ── numeric ascending sort of the result [caller line ~459]
└─ AddCollectivePhysicalCoreIndices(hlo, span) ── write physical_core_indices [caller line ~521]
(backend_config_util::AddCollectivePhysicalCoreIndices, 0x1C868500, absl::Span<long const>)
```text
### 算法
```text
allowed = GetAllowedCores(hlo) // btree_set<long> candidate mask
selected = SelectCores(hlo, allowed, devcount, cost, // 5-phase greedy filter; UNSORTED
assigned, reach, assign_groups)
sort(selected, ascending) // numeric — in the CALLER, not SelectCores
AddCollectivePhysicalCoreIndices(hlo, span(selected)) // long -> int32, write proto fieldGOTCHA —
SelectCores返回的 vector 处于{phase, cost}顺序，不是数值顺序。生成最终physical_core_indices的数值__sort位于调用方（0x10FDF480），在SelectCores返回之后。若重新实现者在选择例程内部排序，或把选择顺序当成 proto 顺序，当列表随后被截断到 device/megachip 数量时会错误地给核心排名 —— 因为调用方排序后截断会丢弃编号最高的核心，但决定哪些核心会被追加的是phase-and-cost 顺序。
相关函数
| Function | Address | 作用 |
|---|---|---|
SparseCoreQueueAssignment::GetAllowedCores | 0x10FDA3C0 | 候选核心 mask（btree_set） |
SparseCoreQueueAssignment::SelectCores | 0x10FDC4E0 | 五阶段有序选择 |
SparseCoreQueueAssignment::AssignQueueIDsToAsyncStart | 0x10FDF480 | 逐 collective driver；调用二者，然后排序 |
backend_config_util::AddCollectivePhysicalCoreIndices | 0x1C868500 | 从 Span<long const> 写入 physical_core_indices |
GetAllowedCores — 候选核心 Mask
目的
GetAllowedCores（0x10FDA3C0）是一个成员函数（sret），签名为 GetAllowedCores(HloInstruction*) → btree_set<long,less<long>,allocator<long>,256>。它返回 SelectCores 可为一个 collective 选择的物理 SC core ID 池。该池是 collective 的 megachip parallelism 覆盖的 chip 与该 collective scheduling-resource reservation 下仍可用核心的交集。
入口点
GetAllowedCores (0x10FDA3C0) ── sret = btree_set<long>; rdx = hlo
├─ GetChipIDsFromParallelismConfig(hlo) (0x10FDBF40) ── vector<long> megachip per-axis chip IDs
│ └─ GetMegaChipParallelism (0x1C867B00) ── StatusOr<InlinedVector<long,4>>; per-axis split
├─ walk device-assignment member set ── Swiss-table SOO: control [hlo+0x90], slots [hlo+0x98], slot stride 0x68 (13 qwords)
│ └─ GetSparseCoreResources(member-op) (0x10FDC0A0)── btree_set<long> resource-type IDs {0,23..28}
├─ flat_hash_map<resource_id, btree_set<chip_id>> ── @ policy global 0x2181D940 (group chips by resource)
├─ flat_hash_map<chip_id, refcount> ── @ policy global 0x21639C10 (per-chip occupancy)
├─ per-resource thread-local budget ── __tls_get_addr(&qword_22048D78); dec; gate `>= 2`
└─ second pass over [hlo+0xC0] ── re-apply budget; assemble output btree_set
```text
### 算法 — 两个输入
**`GetChipIDsFromParallelismConfig`**（`0x10FDBF40`）读取 `GetMegaChipParallelism`（`0x1C867B00`）—— 也就是 queue-assignment driver 使用的同一个 proto reader —— 并把每个逐轴 `MegaChipParallelism` 值复制进 `RepeatedField<long>`。空 `StatusOr` 是硬错误。因此这里的 "chip IDs" 是逐轴 megachip split 列表，不是 device 列表。
**`GetSparseCoreResources`**（`0x10FDC0A0`）返回 collective 占用的 AsyncTracker scheduling resource-type ID 的 `btree_set<long>`，逐字节确认如下：
```c
function GetSparseCoreResources(op): // 0x10FDC0A0
if op.opcode == 0x11 /* custom-call */:
cfg = GetSparseCoreConfig(op) // 0x1C868D20 (backend_config_util)
if (cfg.has_type() /* type hasbit */): // byte-exact: `(cfg_byte[0x10] & 4) != 0`
switch (collective_type_enum - 1) of {0..6}: // `dec eax; cmp 6; ja default`, jump table @0xAC0A82C
case 0 (enum 1): insert 28 (.rodata 0xAC0A910)
case 1 (enum 2): insert 23 (.rodata 0xAC0A8E8)
case 2 (enum 3): insert 24 (.rodata 0xAC0A8F0)
case 3 (enum 4): // NOT a constant — derive from the wrapped op:
body = op.async_wrapped_instruction() // 0x1E5AA300
.called_computations() // 0x1E5885A0 (root instruction)
insert AsyncTracker::GetResourceTypeForOp(root.opcode) // 0x13612240
case 4 (enum 5): insert 25 (.rodata 0xAC0A8F8)
case 5 (enum 6): insert 26 (.rodata 0xAC0A900)
case 6 (enum 7): insert 27 (.rodata 0xAC0A908)
default: insert 0 (.rodata 0xAC0A920)
// else (type hasbit clear): insert 0 (.rodata 0xAC0A920)
else: // not a custom-call
insert 0 (.rodata 0xAC0A8E0) // single resource-type 0
return result // btree_set<long>, insert_hint_uniqueResource-type long 常量直接从 .rodata 读取（VMA == file offset）：0xAC0A8E0 = 0，0xAC0A8E8 = 23，0xAC0A8F0 = 24，0xAC0A8F8 = 25，0xAC0A900 = 26，0xAC0A908 = 27，0xAC0A910 = 28，以及全零 slot 0xAC0A918 = 0 / 0xAC0A920 = 0（default / hasbit-clear 分支）。每个非零值都是一个不同的 AsyncTracker scheduling resource；每个 budget unit 中一个核心可为每个 resource 持有一个 collective，这正是限制候选集合的机制。注意结构：六个 switch 分支（enum 1,2,3,5,6,7）插入固定常量 {28,23,24,26,27} ∪ {25}；enum-4 分支不插入常量 —— 它 unwrap async op（async_wrapped_instruction → called_computations root），并插入 AsyncTracker::GetResourceTypeForOp(root.opcode) 返回的值。非 custom-call 路径和 default / hasbit-clear 路径插入单个 resource-type 0。
算法 — mask 构建和 reservation budget
主体遍历 device-assignment / replica-group 结构，并通过两个 Swiss table 累积候选池：
function GetAllowedCores(hlo): // 0x10FDA3C0
chips = GetChipIDsFromParallelismConfig(hlo) // megachip per-axis chip IDs
by_res = flat_hash_map<long, btree_set<long>>{} // resource_id -> chip_set (policy @0x2181D940)
occupancy = flat_hash_map<long, long>{} // chip_id -> refcount (policy @0x21639C10)
budget = __tls_get_addr(&qword_22048D78) // &(per-resource thread-local reservation budget)
for assignment in device-assignment (Swiss-table SOO @[hlo+0x90]/[hlo+0x98], slot stride 0x68):
resources = GetSparseCoreResources(member-op) // {0,23..28}
for (res, chip) in (resources × chips):
v = (*budget)-- // post-decrement the in-place TLS long
if v >= 2: // gate (pre-decrement value): chip still reservable
by_res[res].insert(chip) // insert_hint_unique
occupancy[chip] += 1 // inc [slot+8]
// second pass re-applies the budget over [hlo+0xC0]
result = btree_set<long>{}
for chip surviving the budget across by_res / [hlo+0x70] device count:
result.insert(chip) // internal_emplace / insert_hint_unique
return result
```text
> **NOTE —** budget gate 是两个位置上的 **post-decrement-then-compare** 模式 `v = (*budget)--; if (v >= 2)`（每个 map-fill pass 一个；TLS `long` 通过 `__tls_get_addr(&qword_22048D78)` 加载，后递减，然后测试 `>= 2`）。**写入并初始化**该线程本地 `long` 的 writer —— 因而它的初始值 —— 在此函数中不可见。Reservation 是 embedding-device / reserved-core knob：这是把为 embedding device 保留的核心（`sc_dev − num_embedding_devices` split）从候选池中排除的机制；seed 是 `NumEmbeddingDevices`、固定的逐 resource 核心上限，还是 `LogicalDevicesPerChip(SC)`，这里无法确定。该 budget 所属的 resource→limit map 见 [SC Queue Assignment & Reservation](sc-queue-assignment-reservation.md)。
>
> **GOTCHA —** 两个 map global 是*策略* singleton，不是逐调用状态：`flat_hash_map<long, btree_set<long,…,256>>::GetPolicyFunctions()::value` 位于 `0x2181D940`，`flat_hash_map<long, long>` 位于 `0x21639C10`。btree empty-node sentinel 是 `0x2181D930`。重新实现者实例化这些 map 时，必须匹配 resource→chip-set value type 的 `btree_set<...,256>` node fan-out（256）；`256` 是类型的一部分，btree-node walk stride 依赖它。
### 函数图
| Function | Address | 作用 |
|---|---|---|
| `SparseCoreQueueAssignment::GetAllowedCores` | `0x10FDA3C0` | 候选 mask 构建（sret btree_set） |
| `(anon)::GetChipIDsFromParallelismConfig` | `0x10FDBF40` | megachip 逐轴 chip ID → `vector<long>` |
| `(anon)::GetSparseCoreResources` | `0x10FDC0A0` | 逐 collective resource-type ID `{0,23..28}` |
| `GetMegaChipParallelism` | `0x1C867B00` | `StatusOr<InlinedVector<long,4>>` 逐轴 split |
| `backend_config_util::GetSparseCoreConfig` | `0x1C868D20` | offload-config 读取（type enum + hasbit） |
| `AsyncTracker::GetResourceTypeForOp` | `0x13612240` | enum-4 switch 分支：从 unwrapped async-op root opcode 得到 resource-type |
| reservation budget | `__tls_get_addr(&qword_22048D78)` | 逐 resource 排除计数器 |
---
## `SelectCores` — 五阶段贪婪过滤器
### 目的
`SelectCores`（`0x10FDC4E0`）是一个成员函数（sret），返回 `StatusOr<vector<long>>`。它接收 allowed btree-set、device count、一个 `double cost`、已分配 collective 的运行中 `vector<Info>`、`HloReachabilityMap` 和 assignment-group vector，并按贪婪优先级顺序产生核心。符号表中的 demangled 签名：
```text
SelectCores(HloInstruction const* hlo,
btree_set<long, less<long>, allocator<long>, 256> const& allowed,
long devcount, double cost,
vector<Info>& assigned,
HloReachabilityMap& reach,
vector<vector<HloInstruction*>>& assign_groups)
-> StatusOr<vector<long>>Info Struct（sizeof 0x60）
assigned vector 的每个元素描述一个已经放置的 async collective。该 struct 为 0x60 字节（元素 stride 计算为 lea rax,[rax+rax*2]; shl rax,5，即 ×3 ×32 = ×96）。字节布局由逐阶段字段比较确认：
[Info+0x00] HloInstruction* the already-assigned async collective op
[Info+0x08] vector<long> tag low bit (bit0) = is-heap flag; if heap, size = tag>>1
[Info+0x10] vector<long> data the physical cores this assigned op holds (inline if <=4, else heap ptr)
[Info+0x30] NDPlaneInfo the SC ND-plane descriptor:
+0x30 int32 plane#1 | +0x34 int32 #2 | +0x38 int32 #3
+0x3C int32 size_x | +0x40 bool has_x
+0x44 int32 size_y | +0x48 bool has_y
+0x4C int32 size_z | +0x50 bool has_z
+0x54 bool (NDPlaneStrideInfo trailing byte)
+0x58 bool across_cores_on_chip (megacore / both-cores flag)
```text
`NDPlaneInfo` 子布局与 SC collective 栈其他位置使用的 descriptor 匹配；`NDPlaneInfo::ToString`（`0x10FDF2A0`）是 VLOG 调用的 formatter（不存在时用字面量 `"None"`）。
### 算法 — Phase 0 setup 和 cost tie-break
```c
function SelectCores(hlo, allowed, devcount, cost, assigned, reach, assign_groups): // 0x10FDC4E0
// a) count allowed btree elements (leaf-node walk, 0x100 stride)
// b) materialize allowed IDs into a heap vector<long> (allowed_vec)
allowed_vec = materialize(allowed)
// c) the ONLY use of `cost`: order candidates by ascending per-core weight
stable_sort(allowed_vec, $_0) // $_0(a,b): vucomisd weight[a], weight[b] (0x10FE8AE0 / cmp @0x10FE8B34)
// d) get THIS collective's ND plane (the target plane the predicates test against)
target_plane = TryGetNDPlaneInfoForSparseCoreCollectives(hlo, this->target_) // 0x10FDEDC0
if !target_plane.ok(): return Error(src line 269) // AddSourceLocationImpl(..., 269, "…sparse_core_queue_assignment.cc")
selected = []
for phase in [P1, P2, P3, P4, P5]:
for core in allowed_vec: // ascending cost
if core in selected: continue
if phase.predicate(core, assigned, reach, assign_groups, target_plane):
selected.append(core) // grow via _Znwm + memcpy
VLOG(phase.msg)
return Ok(selected) // sret: [+0]=1 ok-tag, [+8]=data, [+0x10]=size, [+0x18]=capQUIRK —
double cost参数只在一个位置被消费：stable_sortcomparator$_0（对逐核心doubleweight array 执行vucomisd）。它只是 tie-break。在同等满足某个 phase predicate 的候选中，低 cost 核心先追加。没有决定核心是否被选中的逐候选 score —— 只有 phase predicate 决定。填充每个核心 weight 的算术（调用方中的 LatencyEstimator / queue-occupancy feed）在此函数中不可见；这里确定的只是 comparator 的使用方式（升序排序）。
算法 — 五个 predicate
每个 pass 都扫描 assigned；SIMD vpcmpeqq（4-wide）搜索会测试候选核心是否出现在某个 Info 的物理核心列表 [Info+0x10..] 中，然后应用逐 phase predicate。VLOG 字符串已在反编译中按源码顺序逐字节确认。
P1 SAME ND PLANE (lambda $_1, src line 279):
exists Info holding `core` whose NDPlaneInfo[+0x30] == target_plane (all 3 ints, the
per-axis has/size optionals, and across_cores[+0x58]) -> append.
VLOG: "Adding core <c> to selected cores for <hlo> because it is running <NDPlaneInfo|None>
on the same ND plane."
P2 DATA DEPENDENCY (lambda $_2, src line 302):
probe HloReachabilityMap `reach`: crc32 Swiss find keyed on (GetModule()->[module+0xAC4],
HloInstruction::unique_id()) for BOTH the Info's op and the target hlo; on hit, test the
reachability bit (row = bitidx>>0xA, bound [reach+0x40]; col = bitidx&0x3FF scaled by
[reach+0x28]; base [reach+0x38]; `bt` test). Reachable in either direction -> append.
VLOG: "Adding core <c> ... because an assigned op <op> has data dependency with it."
P3 ASSIGNMENT-GROUP HINT (lambda $_3, src line 329):
for each group in assign_groups containing `hlo`: for each group member an Info holds and
that appears in allowed_vec -> append that core (this is a PIN).
VLOG: "Adding core <c> ... due to hint from pre-determined assignment groups."
P4 NOT A DIFFERENT PLANE (lambda $_5 add, src line 362 / lambda $_4 skip, src line 352):
same field compares as P1 but the success branch is the negation (`setne; and across_cores`).
candidate NOT found running a collective on a DIFFERENT ND plane -> append.
VLOG (add): "Adding core <c> ... because it is not running a collective on a different ND plane."
VLOG (skip): "Core <op> is running <NDPlaneInfo|None> on a different ND plane."
P5 FALLBACK FILL (no VLOG):
every candidate in allowed_vec not already selected -> append unconditionally.
```text
| Phase | Src line | Predicate（候选核心 `c`，扫描 `assigned`） | VLOG |
|---|---|---|---|
| P1 | 279 | ∃ `Info` 持有 `c`，且其 `NDPlaneInfo == target`（**same plane**） | "… because it is running … on the same ND plane." |
| P2 | 302 | ∃ `Info`，其 op 与 `hlo` 可达 ↔（**data dependency**） | "… has data dependency with it." |
| P3 | 329 | `hlo` ∈ 某个 assignment group ∧ 某个成员的 `Info` 持有 `c` | "… due to hint from pre-determined assignment groups." |
| P4 | 362 | **没有**持有 `c` 的 `Info` 位于**不同** ND plane | "… because it is not running a collective on a different ND plane." |
| P4 (skip) | 352 |（`c` **在**不同 plane 上 → skip） | "Core … is running … on a different ND plane." |
| P5 | — | always（追加每个剩余 allowed core） |（无） |
> **NOTE —** `SelectCores` 内部的 assignment-group VLOG 是**正向** pin（"Adding core … due to hint …"，第 329 行）。互补诊断 "Not pinning … due to hint from pre-determined assignment groups." **没有**在 `SelectCores` 或 `GetAllowedCores` 内部被引用；它位于更高层的 assignment-group 决策（group construction / `AssignQueueIDsForComputation` 层）。`SelectCores` 自己的逐候选 "skip" 诊断是 P4 的 "Core … is running … on a different ND plane."（第 352 行）。
### 返回形状
`SelectCores` 把 `StatusOr<vector<long>>` 写入其 sret：`[sret+0] = 1`（ok tag），`[sret+8] = data ptr`，`[sret+0x10] = size`，`[sret+0x18] = capacity`。Vector **未排序** —— 其顺序精确为 `{P1 same-plane} ++ {P2 data-dep} ++ {P3 assign-group} ++ {P4 not-different-plane} ++ {P5 all-remaining}`，每个 pass 内 cost 升序。调用方在把它复制（截断 `long → int32`）到 `physical_core_indices` 前，会对它进行数值排序。
> **GOTCHA —** `SelectCores` 中**没有**闭式逐候选数值 score，也**没有**显式 core→queue 双射算术。Placement 是确定性贪婪优先级：候选的位置为 `(phase index 1..5, then ascending per-core cost)`，其中 phase 成员由持有该核心的已分配 async collective 是否 same-plane / data-dependent / 在同一 assignment group / 在不同 plane 上决定。重新实现者必须复现*pass 顺序*和*cost tie-break*，而不是发明标量 score。
### 函数图
| Function | Address | 作用 |
|---|---|---|
| `SparseCoreQueueAssignment::SelectCores` | `0x10FDC4E0` | 五阶段有序选择（sret StatusOr） |
| `…::SelectCores(...)::$_0`（在 `__stable_sort` 中） | `0x10FE8AE0` | cost 升序 comparator（`vucomisd` @ `0x10FE8B34`） |
| `(anon)::TryGetNDPlaneInfoForSparseCoreCollectives` | `0x10FDEDC0` | predicate 使用的目标 ND-plane |
| `NDPlaneInfo::ToString` | `0x10FDF2A0` | VLOG plane formatter（或字面量 `"None"`） |
| `HloReachabilityMap` probe | — | P2 bit-matrix `bt` test（crc32 Swiss find） |
| `HloInstruction::async_wrapped_instruction` | `0x1E5AA300` | 用于 resource 派生的 unwrap（与 mask 共享） |
---
## `tensor_split_factor` 交互
### 目的
相关字段是 collective offload-config 的 **`tensor_split_factor`**（此 build 中没有 `tensor_split_mode` 字段 —— “mode” 由 `tensor_split_factor() > 1` 是否成立隐含）。它**不会**出现在 `SelectCores` 或 `GetAllowedCores` 内部。相反，它在 collective ring-strategy 层的*下游*被消费，控制*所选*核心如何分区一个 collective 的 tensor。因此它与核心选择的关系是数量约束，而不是选择输入。
### 它位于哪里
字符串 `tensor_split` 和 `tensor_split_factor` 解析到 SC collective strategy 代码，而不是 queue-assignment pass。两个逐字节确认的位置：
```text
SinglePhaseAGStrategy::AdjustStrategiesForSingleCoreIfNeeded (0x1338C1E0):
gate: "offload_config().use_single_sparse_core()
|| (offload_config().has_tensor_split_factor()
&& offload_config().tensor_split_factor() > 1)"
error: "Ring adjustment is only supported either for single core or tensor-split mode."
emitter_helpers::CreateRingStrategiesForNdFromExplicitTable (0x13390900):
RET_CHECK: "ici_strategy_config.color_strategies_size() % tensor_split_factor.value_or(1) == 0"
(offload_collective_strategies.cc:3966)因此 mode 是：single-core（use_single_sparse_core() → collective 折叠为一个 SC core）与 tensor-split（tensor_split_factor() > 1 → tensor 跨那么多个核心拆分）。Ring strategy 要求 color strategy 数量能被 tensor_split_factor 整除（默认 1）。
它如何关联选择
SelectCores/GetAllowedCores 产生一个 collective 可用的物理核心集合；tensor_split_factor 决定单个 collective 的 ring 实际扇出到其中多少个核心，整除性 RET_CHECK 确保 strategy table 能均匀分区。选择策略对 split mode 无感 —— 它既不读取 tensor_split_factor，也不根据 use_single_sparse_core 分支。重新实现者应把 tensor_split_factor 当作发出的 strategy 的属性处理，它必须与所选核心数量一致，并通过 offload-config 传播；该 offload-config 由 GetSparseCoreResources 用于 resource-type enum 的同一个 GetSparseCoreConfig reader 读取。
NOTE —
tensor_split_factor↔ selected-core-count 一致性的 gate 字符串和整除性RET_CHECK都位于消费physical_core_indices的 collective strategy 层。tensor_split_factor为 N 时映射到所选核心列表中哪 N 个具体条目的精确算术，属于 ring-strategy 构造的一部分，超出本选择页面范围（它属于 offload-strategy emitter）。tensor_split_factor是所选核心集合的 consumer，不是选择输入。
函数图
| Function | Address | 作用 |
|---|---|---|
SinglePhaseAGStrategy::AdjustStrategiesForSingleCoreIfNeeded | 0x1338C1E0 | single-core vs tensor-split gate |
SinglePhaseRSStrategy::AdjustStrategiesForSingleCoreIfNeeded | 0x1338A7A0 | reduce-scatter analogue |
emitter_helpers::CreateRingStrategiesForNdFromExplicitTable | 0x13390900 | color_strategies_size() % tensor_split_factor == 0 |
backend_config_util::GetSparseCoreConfig | 0x1C868D20 | offload-config reader（与 mask 共享） |
逐代际说明
SelectCores 或 GetAllowedCores 的代码中没有按代际分支。候选池随 chip 的 megachip parallelism 和 SC count 变化（通过 GetMegaChipParallelism 和 device assignment 数据驱动），Info[+0x58] 中的 across_cores_on_chip flag 是 plane predicate 读取的 megacore / both-cores 判别符 —— 但策略（五个 pass、cost tie-break、resource-reservation 交集）在所有硅片上相同。约束一个 chip 暴露多少物理 SC core 的 4:1 SparseCore:TensorCore 比例在 SparseCore Hardware Architecture（SparseCoreCountPerTensorCore，0x1C6CB760）中确立；本页消费该数量，而不计算它。
相关组件
| Name | 关系 |
|---|---|
AssignQueueIDsToAsyncStart（0x10FDF480） | 调用 GetAllowedCores + SelectCores，然后排序结果的 driver |
AddCollectivePhysicalCoreIndices（0x1C868500） | 把已排序 core list 写入 physical_core_indices 的 sink |
GetSparseCoreConfig（0x1C868D20） | resource-type enum 和 tensor_split_factor 共享的 offload-config reader |
AsyncTracker::GetResourceTypeForOp（0x13612240） | enum-4 switch 分支的 resource-type 派生（来自 unwrapped async-op root opcode），供给 GetSparseCoreResources |
交叉引用
- SparseCore 概览 — Part IX 的导航入口；engine 名称、逐代际存在性、data path。
- SparseCore 硬件架构 — 几何来源（
SparseCoreTarget）、4:1 SC:TC 比例，以及此策略选择的物理核心数量。 - SC Backend Pipeline — queue-assignment pass 所在的 SC-MLO pass 流水线（以及 MEGACORE barrier）。
- SC Queue Assignment & Reservation —
GetAllowedCores逐 resource budget 所属的 resource→limit reservation map。 - Physical-Core Placement —
physical_core_indices的 collective 侧 consumer，所选 core list 在此映射到 ring/megachip topology。 - GetSparseCoreConfig — 读取 resource-type enum 和
tensor_split_factor的 offload op-type 配置。 - OneSlot Router — 所选核心供给的逐 slot routing。
- getSequencerType — SCS/TAC/TEC engine-selection 函数和 sequencer-type enum。
- Region → Sequencer Outliner — 把所选 SC computation 分区为逐 engine bundle stream。
- SCS（Scalar）Engine · TAC Engine · TEC（Vector）Engine — 所选核心运行的三个 sub-engine bundle surface。
- Binary:
extracted/libtpu-0.0.40-cp314-cp314-manylinux_2_31_x86_64/libtpu/libtpu.so（build-id89edbbe81c5b328a958fe628a9f2207d） - 索引条目: Part IX — SparseCore & BarnaCore / SparseCore back-end — 返回索引