SC 内核选择和卸载门

此页上的每个地址、偏移量、字段号和常量都是从 libtpu-0.0.40-cp314 轮中的 libtpu.so 中精确读取的字节(BuildID md5 89edbbe81c5b328a958fe628a9f2207d，构建 libtpu_lts_20260413_b_RC00)。 .text VMA == 文件偏移量(0xE63C000); .rodata VMA == 文件偏移量(0x84A0000); .data.rel.ro VMA - 0x200000 == 文件偏移量。其他版本有所不同。

摘要

此页面是 SparseCore (SC) 核心选择的集体堆栈视图以及打开 SC 卸载基板的 卸载门。它回答了集合堆栈的重新实现者在构建任何 SC 环配置之前必须回答的两个问题：(1) 该集合是否已路由到 SC 基板， 和 (2) 一旦路由，它在哪个物理 SC 内核上运行，以什么顺序运行？ 这两个问题都在一个后端通道 xla::jellyfish::SparseCoreQueueAssignment 内决定，并且两个门都在相同的后端配置枚举：SparseCoreConfig.offload (字段 2，输入 xla::jellyfish::Offload)，通读 backend_config_util::GetSparseCoreConfig (0x1C868D20)。

核心选择策略是一个双功能生产者。 GetAllowedCores (0x10FDA3C0) 计算候选核心掩码 — 按集体调度资源类型 ID {0, 23..28} 分组的兆芯片每轴芯片 ID，减去每资源预留预算已耗尽的核心。 SelectCores (0x10FDC4E0) 通过在成本排序的候选者上运行五阶段贪婪优先级过滤器，将该掩码转换为有序 vector<long>，其中每个阶段扫描已分配的异步集合的正在运行的 vector<Info>。阶段顺序 — 相同 ND 平面 → 数据依赖 → 分配组提示 → 不在不同平面上 → 回退 — 是放置策略；没有数字记分器。此页面记录了集体堆栈使用该策略的级别，以及决定该策略是否运行的门。两个函数(Info 结构、SIMD 核心搜索、HloReachabilityMap 位探测、瑞士表掩码构建)的 字节级主体 均归 SC内核选择 所有 - 此页面链接它，它不会重新派生它。

该页面由三个单元组成：卸载门(RunHloScheduler的能力/平台谓词和选择资源臂的SparseCoreConfig.offload op类型分类器)、GetAllowedCores(从集体的角度来看，候选掩码和预留预算)和**SelectCores**(五阶段政策和成本决胜局)。数字 __sort 到 physical_core_indices 以及每颜色→物理核心索引映射位于下游并位于 物理核心放置 上。

重新实现，合约为：

一个枚举控制路由和资源选择。 SparseCoreConfig.offload(字段 2)决定集合体是否为 SC 卸载(OFFLOAD_COLLECTIVE = 4 到达卸载集合体配置构建器)，并且在 GetSparseCoreResources 中，集合体占用哪个调度资源类型 ID {0, 23..28}(这反过来又决定集合体占用的调度资源类型 ID)。候选掩码)。
硬件门是功能 has-bit 或平台 bool。 (Target[+0x628] & 4) ∨ (Target[+0x540] ≠ 0) — SC 卸载功能 has-bit(真实硬件)或 iss/模拟器平台 bool — 均以 Target::Init 编写，以 TpuVersion == 5 为每代基础。
核心选择是候选掩码，然后是有序过滤器，而不是评分器。 GetAllowedCores → btree_set<long> 掩码； SelectCores → 有序 vector<long>(返回未排序 — 调用者按数字排序)。核心的等级是(phase 1..5) Major，(ascending per-core cost) Minor。
预留预算是预留核心排除。 每个资源线程本地 long (__tls_get_addr(&qword_22048D78))，每次分配递减并门控 >= 2，删除对该资源的预留已耗尽的核心 - 嵌入设备预留核心效应。


卸载门	`SparseCoreCompiler::RunHloScheduler` (`0x1306F820`) — 5 项谓词
能力位	`Target[+0x628] & 4`(SC 卸载 has 位)∨ `Target[+0x540]`(平台 == 2) — 在 `Target::Init` (`0x1D60FC20`) 中设置
每代基础	`TpuVersion == 5` (`TpuChipParts[+0] == 5`)
Op 类型枚举	`SparseCoreConfig.offload`(字段2，`xla::jellyfish::Offload`)；结构体 `+0x24`，有位 `+0x10` 掩码 `0x4`
配置读取器	`backend_config_util::GetSparseCoreConfig` (`0x1C868D20`)
候选掩码	`GetAllowedCores(HloInstruction*)` (`0x10FDA3C0`) → `btree_set<long,…,256>`
订购选择	`SelectCores(hlo, allowed, devcount, cost, assigned, reach, assign_groups)` (`0x10FDC4E0`) → `StatusOr<vector<long>>`
资源类型集	`{0, 23, 24, 25, 26, 27, 28}`(`GetSparseCoreResources`、`0x10FDC0A0`)
预订预算	每个资源线程本地 `long` (`__tls_get_addr(&qword_22048D78)`)，递减，门控 `>= 2`
选择阶段	P1 同平面 · P2 数据深度 · P3 分配组 · P4 非不同平面 · P5 后备
抢七	成本提升型`stable_sort`比较器`$_0`(`vucomisd`、`0x10FE8AE0`)
信心	已确认(反编译 + 二进制字节锚定)，除非行或标注另有说明

1. 卸载门

用途

在选择任何 SC 内核之前，编译必须决定 SparseCore 卸载延迟隐藏调度程序是否运行。该决定是硬件/平台功能检测(这部分是否支持 SC 卸载，还是模拟器？)与模块内容检查(HLO 模块实际上包含 SC 指令吗？)和标志相结合。大门位于SparseCoreCompiler::RunHloScheduler (0x1306F820)；当它不成立时，集体停留在密集的 TensorCore 路径上，并且 GetAllowedCores / SelectCores 都不会运行。

该页面拥有门位——每个术语的含义及其写入位置。门在基板分割中的角色(哪个调度程序安装哪个异步跟踪器)位于**On-Pod 集体 — 剖面图**；不要在这里重复该框架。

算法

完整谓词，在 RunHloScheduler 反编译中进行字节确认(偏移量 1576 = 0x628、1344 = 0x540、148 = 0x94)：

function RunHloScheduler(target, module):                 // 0x1306F820
    chip_cfg = target->chip_config;                        // target[+119*8][+24]
    runSC =  TpuChipConfig::Megachip(chip_cfg)             // line 80 / 193
          && CoresPerChip(kSparseCore) > 0                 // chip_cfg[+0x94] > 0
          && ( (target[+0x628] & 4) != 0                   // SC-offload-capability has-bit
               || target[+0x540] != 0 )                    // platform == 2 (iss/simulator)
    if !runSC: return;                                     // stay on dense TC path
    if !ModuleContainsLEMSparseCoreInstruction(module):    // line 196 (0x13853280)
        return;
    if !ReadOneBool(FLAGS_xla_sc_enable_latency_hiding_scheduler):  // line 243
        return;
    ... run the SparseCore latency-hiding schedulers ...
```text

反编译条件逐字：

```text
*(int *)(v8 + 148) > 0
  && ((*((_BYTE *)this + 1576) & 4) != 0 || *((_BYTE *)this + 1344))     // line 85

并在模块内容检查之前在 SC 路径内的第 193 行处进行相同的重播。

两个目标位

两个 Target 字段均写入 jellyfish::Target::Init (0x1D60FC20)：

字段	含义	设置在哪里	如何
`Target[+0x540]`(字节/布尔)	平台类型==2(`iss`/模拟器平台)	`Target::Init` `@0x1D610B1B`	`sete` 来自 `(TpuTopology[+0] == 2)`
`Target[+0x541]`(字节/布尔)	平台类型 == 1 (`grm`)	`Target::Init` `@0x1D610B29`	`sete` 来自 `(TpuTopology[+0] == 1)`
`Target[+0x628]` 位 0 (`& 1`)	配置子字段有位(同伴)	`Target::Init` `@0x1D611D52`	`\|= 0x1`
`Target[+0x628]` 位 2 (`& 4`)	SC 卸载能力有位	`Target::Init` `@0x1D612121`	`\|= 0x4`，位于由 SC 卸载功能检测门控的配置附加循环内

在实际硬件上，门由位 2 承载(谓词门控配置附加循环内的每代功能集)。无论功能位如何，platform == 2 分支都会强制采用 iss 模拟器的 SC 路径。相同的特征检测谓词 - Megachip ∧ CoresPerChip(SC)>0 ∧ ((+0x628&4) ∨ +0x540) ∧ GetContinuationQueues(SC)[0]==2 - 在 Target::Init 本身内部重播两次(@0x1D611A4C、@0x1D611E52)。

注意 — TpuTopology[+0] 是拓扑的第一个标量 = 平台类型枚举(内部 {0 hardware, 1 grm, 2 iss}，描述符字符串顺序)。闸门采用== 2；平台枚举→名称配对是描述符顺序，不单独进行开关确认，但无论如何，门角色都是字节精确的。 (对枚举器排序的信心较低；在 == 2 比较中已确认。)

每代硬件基础

SC 卸载并发旋钮在一代芯片中默认打开。 ShouldEnableConcurrentSparseCoreOffloading (0x1D6B6F80) 和 EnableSparseCoreOffloadQueuingInLhs (0x1D6B81E0) 都将其硬件默认值计算为 (TpuChipParts[+0] == 5)，即 TpuVersion == 5，然后让 AutoOr<bool> 原型标志覆盖 (test eax, 0x100; cmove)。 TpuChipParts[+0] 是基于 0 的内部芯片生成枚举。

`TpuVersion`(内部)	代号	原始值(= 内部 + 1)
0	水母	1
1	龙鱼	2
2	河豚	3
3	毒蛇鱼	4
4	鬼石	5
5	`6acc60406`(在此版本中混淆的代号)	6 — SC 卸载默认开启

奇特 — 大门有两个出口匝道，看起来很像，但其实不然。 Target[+0x628]&4 是在 Target::Init 发现的每代功能； FLAGS_xla_sc_enable_latency_hiding_scheduler 是一个运行时标志； ShouldEnableConcurrentSparseCoreOffloading 是第三个 AutoOr<bool> 旋钮，默认为 TpuVersion==5。将“功能位”与“启用标志”混为一谈的重新实现将误选通模拟器(其中 +0x540 强制启用)，并误选通标志强制打开的较旧部件(其中位 2 已清除，但选通仍然失败)。

2. `SparseCoreConfig.offload` — op 类型分类器

用途

一旦门保持，SparseCore 操作类型将从一个后端配置枚举中读取：SparseCoreConfig 字段 2 offload，类型为 .xla.jellyfish.Offload 的 TYPE_ENUM。它不是自定义调用目标名称，也不是 MLIR 操作类型。候选掩码生产者(GetSparseCoreResources，本页§3)和密集端调度程序的资源生产者(MayAddSparseCoreResource)都读取相同字段。该枚举将每个集合路由到一个调度资源类型 ID，该 ID 是预留预算计量和掩码组下的每个集合的资源。

枚举所在的位置

GetSparseCoreConfig (0x1C868D20) 返回完整的 SparseCoreConfig 原型；消费者直接从返回的结构中读取卸载字段。在 GetSparseCoreResources 中，读取是字节确认的：

if ( *((_BYTE *)hlo + 12) == 17 )            // opcode == 0x11 (custom-call)        line 44
{
    GetSparseCoreConfig(&cfg, hlo);          //                                     line 46
    if ( (cfg[16] & 4) != 0 )                // has-bit [cfg+0x10] mask 0x4 = field 2 line 47
    {
        switch ( offload_enum )              // enum at [cfg+0x24], indexed enum-1   line 49
        { case 1: ... case 7: ... }
    }
}
else
    GetResourceTypeForOp(async_body_opcode); //                                     line 142
```text

has 位偏移(`+0x10` 掩码 `0x4`)和值偏移(`+0x24`)与 `SparseCoreConfig::_InternalSerialize` 字段映射(字段 2，has 位 `0x4`，存储在 `+0x24`)匹配。完整的原始字段映射和描述符字节位于 **[SparseCoreConfig(GetSparseCoreConfig)](../sparsecore/getsparsecoreconfig.md)**。

### `Offload` 枚举 → 资源臂

```text
xla::jellyfish::Offload (9 enumerators, EnumDescriptorProto-confirmed):
  0  OFFLOAD_UNSPECIFIED
  1  OFFLOAD_EMBEDDING
  2  OFFLOAD_GATHER
  3  OFFLOAD_SCATTER
  4  OFFLOAD_COLLECTIVE          ← reaches the offload COLLECTIVE config builder
  5  OFFLOAD_DATA_FORMATTING
  6  OFFLOAD_KERNEL
  7  OFFLOAD_SORT
  8  OFFLOAD_COMPUTE

预订地图制作者 GetSparseCoreResources 通过 enum − 1 对交换机进行索引，因此其活动臂是 {1..7}；密集调度器的 MayAddSparseCoreResource 索引为 enum − 2，因此其活动臂为 {2..7}。每个映射都会产生字节确认的：

`Offload` 值	枚举器	`GetSparseCoreResources`(idx = 枚举 - 1)	`MayAddSparseCoreResource`(idx = 枚举 − 2)
0	`OFFLOAD_UNSPECIFIED`	(无臂)	(无臂；`rt22 ×N-cores` 路径)
1	`OFFLOAD_EMBEDDING`	`rt28`(嵌入；`unk_AC0A910` = `0x1C`)	(无臂；`rt22 ×N-cores` 路径)
2	`OFFLOAD_GATHER`	`rt23` `kSparseCoreGather`	`rt23`
3	`OFFLOAD_SCATTER`	`rt24` `kSparseCoreScatter`	`rt24`
4	`OFFLOAD_COLLECTIVE`	集体臂(异步主体递归)	异步体递归
5	`OFFLOAD_DATA_FORMATTING`	`rt25` `kSparseCoreDataFormatting`	`rt25`
6	`OFFLOAD_KERNEL`	`rt26` `kSparseCoreKernel`	`rt26`
7	`OFFLOAD_SORT`	`rt27` `kSparseCoreSort`	`rt27`
8	`OFFLOAD_COMPUTE`	(超出枚举 - 1 范围)	(无臂；`rt22 ×N-cores` 路径)

GetSparseCoreResources 返回集体占用的资源类型 ID 的集合(它遍历的每个设备分配条目一个)。对于非自定义调用操作，它会解开 async_wrapped_instruction 并在异步主体上使用 AsyncTracker::GetResourceTypeForOp (0x13612240)。它返回的集合 - {0, 23, 24, 25, 26, 27, 28} - 是 §3 的掩码将芯片分组的每个集合的资源密钥。

GOTCHA — 两个索引库(enum − 1 用于保留映射，enum − 2 用于密集调度器)不是转录错误；他们是故意的。预留映射还涵盖 OFFLOAD_EMBEDDING(1)(密集调度程序通过 rt22 ×N-cores 后备进行路由)。对两个生产者使用同一个基的重新实现将默默地删除或重复计算嵌入资源。

3. `GetAllowedCores` — 候选核心掩码

用途

GetAllowedCores (0x10FDA3C0) 为一个集体计算“允许”运行的一组物理 SC 核心 — 然后对候选池 SelectCores 进行排序。从集体堆栈的角度来看，这是两个集体级事实相交的地方：巨型芯片并行性(集体的数据如何跨芯片分割)和每个集体的调度资源(预留预算计量)。字节级瑞士表机制位于**SC内核选择**；本节记录了集体堆栈重新实现者必须重现的策略输入和保留交互。

两个输入

function GetAllowedCores(hlo) -> btree_set<long>:          // 0x10FDA3C0
    chips     = GetChipIDsFromParallelismConfig(hlo)       // line 204 (0x10FDBF40)
    // chips = MegaChipParallelism per-axis split (GetMegaChipParallelism 0x1C867B00)
    for each device-assignment entry in hlo:               // [hlo+0x90]/[hlo+0x98], stride 0x68
        resources = GetSparseCoreResources(member_op)      // line 241/725/930 (0x10FDC0A0)
        // resources ⊆ {0,23,24,25,26,27,28}  — the offload-enum arm from §2
        budget    = *__tls_get_addr(&qword_22048D78)       // line 237/747 per-resource long
        if budget >= 2:                                    // line 244/637/811 — reservation gate
            insert chip into resource's chip_set           // insert_hint_unique (line 522)
            budget--                                        // (*addr)-- (line 636/810)
    ... second pass over [hlo+0xc0] re-applies the budget ...
    return assembled btree_set<long> of surviving cores
```text

- **`GetChipIDsFromParallelismConfig` (`0x10FDBF40`)** 读取 `GetMegaChipParallelism` (`0x1C867B00`) — 调用者 `AssignQueueIDsToAsyncStart` 使用相同的 `MegaChipParallelismConfig` 每轴拆分 — 并发出一个 `long`“芯片ID”每个轴值。这是集体的百万芯片数据分割；它是候选掩码的列源。
- **`GetSparseCoreResources` (`0x10FDC0A0`)** 是§2 生产者：集体占用的资源类型 ID `{0, 23..28}`。这是行键。

### 预留预算 — 预留核心排除

这两个输入通过两个 Swiss 表(`flat_hash_map<resource_id, btree_set<chip_id>>` 和 `flat_hash_map<chip_id, refcount>`)进行累积，由 **每资源线程本地预留预算** 控制。预算是通过 `__tls_get_addr(&qword_22048D78)` 达到的 `long`，每个接受的分配 (`(*addr)--`) 递减一次，并且接受路径被门控 `if (budget >= 2)`。一旦资源的预算低于 2，芯片就不再是该资源的候选者——这就是嵌入设备/保留核心效应：为嵌入设备保留的核心(从 `NumEmbeddingDevices` = `sc_dev − num_embedding_devices` 开始)被排除在候选池之外。

> **QUIRK —** 这里的预留预算是形成候选集的*实时*排除，它与 `SparseCoreQueueAssignment` `[this+0xC0]` `btree_map<long,long>` 预留映射 **不同** — 在 v0.0.40 中，它是在 **没有阅读器** 的情况下构建和释放的(请参阅[SC队列分配和预留](../sparsecore/sc-queue-assignment-reservation.md))。线程本地预算实际上门控`GetAllowedCores`； `[this+0xC0]` 地图没有。将 `[this+0xC0]` 映射连接到核心选择的重新实现者将复制残留成员，而不是实时策略。
>
> **注意 —** 预算的 **初始化程序** 无法追溯到其作者。减量 (`(*addr)--`) 和 `>= 2` 门是字节精确的；无论初始值是 `NumEmbeddingDevices`(每个资源的固定核心上限为 2)还是 `LogicalDevicesPerChip(SC)` 都是 **(低置信度)** — 请参阅下面的“我们没有”。

---

## 4. `SelectCores` — 五阶段安置政策

### 用途

`SelectCores` (`0x10FDC4E0`) 采用 `GetAllowedCores` 候选掩码并生成物理 SC 内核的“有序”`vector<long>`。该顺序是放置策略：集体优先与其共享结构的已分配集体位于同一位置(相同的 ND 平面、数据依赖性、分配组)，并且后备保证覆盖范围。 `0x28E0` 字节主体 — `Info` 结构布局、SIMD 每核搜索、`HloReachabilityMap` 位矩阵探测、每相增长/追加 — 由 **[SC内核选择](../sparsecore/sc-core-selection.md)** 所有；本节记录了集体堆栈读取级别的*策略*(五个谓词、它们的顺序和抢七局)。

### 该策略是贪婪优先级过滤器，而不是记分器

`SelectCores` 中没有封闭形式的 `score(core)`，也没有核心→队列双射算法。候选人的最终排名是 `(phase index 1..5)` 大调，`(ascending per-core cost)` 小调：

```c
function SelectCores(hlo, allowed, devcount, cost, assigned, reach, assign_groups):  // 0x10FDC4E0
    allowed_vec  = materialize(allowed)                      // btree_set -> flat vector<long>
    stable_sort(allowed_vec, $_0)                            // ascending per-core cost (0x10FE8AE0)
    target_plane = TryGetNDPlaneInfoForSparseCoreCollectives(hlo, target_)   // line 425
    selected = []
    for phase in [P1 same-plane, P2 data-dep, P3 assign-group, P4 not-diff-plane, P5 fallback]:
        for core in allowed_vec:                             // ascending-cost order
            if core in selected: continue
            if phase.predicate(core, assigned, reach, assign_groups, target_plane):
                selected.append(core)                        // + per-phase VLOG
    return StatusOr(selected)                                // UNSORTED; caller numerically sorts

SelectCores 返回向量未排序；数字 __sort 到 physical_core_indices 发生在调用者 AssignQueueIDsToAsyncStart (0x10FDF480) 中 - 排序和原型写入位于 物理核心放置 上。

五个谓词

每个谓词扫描已分配的异步集合的正在运行的 vector<Info>(Info 结构携带分配的操作、其物理核心列表及其 ND 平面描述符)。下面的 VLOG 字符串都是从反编译中逐字读取的，并固定源行和阶段标识：

相	VLOG线	候选核心 `c` 的谓词(扫描 `assigned`)	VLOG片段(字节读取)
P1 同平面	`0x117` (279)	∃ `Info` 持有 `c`，其 ND 平面 == 目标的 ND 平面	`… because it is running … on the same ND plane.`
P2 数据深度	`0x12E` (302)	∃ `Info` 持有操作可达的 `c` ↔`hlo` (`HloReachabilityMap`)	`… because an assigned op … has data dependency with it.`
P3 分配组	`0x149` (329)	`hlo` ∈ 某分配组 ∧ 组成员的 `Info` 持有 `c`	`… due to hint from pre-determined assignment groups.`
P4非差分平面	`0x16A` (362)	NO `Info` 持有 `c` 位于不同 ND 平面 → 追加	`… because it is not running a collective on a different ND plane.`
P4(跳过)	`0x160` (352)	`c` 在不同平面上 → 跳过	`Core … is running … on a different ND plane.`
P5 后备	—	始终(附加每个剩余的允许核心)	(无)

P5保证所选集合覆盖整个允许集合； P1-P4 仅修复顺序，它在调用者将列表截断为设备/巨型芯片计数后确定成员资格。这些阶段形成了严格的优先级：任何早期阶段声称的核心都不会被后续阶段重新评估(if core in selected: continue 重复检查)。

抢七局

在一个阶段内，allowed_vec 的单个 stable_sort 打破了每核成本的上升(比较器 $_0，每核 double 权重阵列上的 vucomisd，在 __stable_sort 中) 0x10FE8AE0)。 double cost 参数是数字权重输入 SelectCores 的唯一位置，并且在设置时消耗一次以对候选进行排序 - 而不是每个阶段。在同样满足阶段谓词的核心中，首先附加成本较低的核心。

注意 — 每个核心成本权重的 生产者(构建每核心 double 数组和调用者中的 cost 参数的 LatencyEstimator / 队列占用提要)未跟踪。比较器的use(升序排序)是字节确认的；计算每个核心成本的算法是**(低置信度)** - 决胜顺序在结构上是固定的，但不是封闭形式的。请参阅我们没有什么。

5. 验证说明

卸载门、卸载枚举分类器和五阶段策略与 libtpu.so v0.0.40 的 IDA 反编译进行了交叉检查：
SparseCoreCompiler::RunHloScheduler (0x1306F820)：谓词 *(int *)(v8 + 148) > 0 && ((*((_BYTE *)this + 1576) & 4) != 0 || *((_BYTE *)this + 1344))(偏移量 0x94/0x628/0x540)，Megachip 调用， ModuleContainsLEMSparseCoreInstruction 和 FLAGS_xla_sc_enable_latency_hiding_scheduler — 全部存在(反编译行 80/85/193/196/243)。
GetSparseCoreResources (0x10FDC0A0)：操作码 [hlo+12] == 17 (0x11)、GetSparseCoreConfig，有位 (cfg[16] & 4) (+0x10 掩码0x4)、卸载枚举 switch 臂 case 1..7 和非自定义调用 GetResourceTypeForOp 路径 — 准确(第 44/46/47/49/142 行)。
GetAllowedCores(0x10FDA3C0)：GetChipIDsFromParallelismConfig(第204行)，GetSparseCoreResources(第241/725/930行)，__tls_get_addr(&qword_22048D78)预算(第237/747行)，递减(*addr)--(第 636/810 行)、>= 2 门(第 244/637/811 行)和 insert_hint_unique(第 522 行) — 准确。
SelectCores (0x10FDC4E0)：stable_sort 比较器 $_0(第 298/416 行)、TryGetNDPlaneInfoForSparseCoreCollectives(第 425 行)和所有五相 VLOG 片段 (… on the same ND plane. / … has data dependency with it. / … due to hint from pre-determined assignment groups. / … not running a collective on a different ND plane. / … on a different ND plane.) — 准确。
[低] 三个剩余未知数，每个都在上面标记：(1)每核成本权重 生产者(抢七顺序在结构上是固定的，不是封闭形式的)； (2) GetAllowedCores预留预算初始化器(递减量和>= 2门是精确的；初始值为NumEmbeddingDevices/fixed-2/LDPC(SC)之一)； (3) Target[+0x540] 的 == 2 后面的平台型枚举器排序(比较准确；{hardware,grm,iss} 枚举器顺序是描述符字符串顺序，而不是开关确认)。

我们还没有什么

每核心成本权重生成器。 stable_sort 比较器 $_0 按升序 cost[core] 对 allowed_vec 进行排序，由 double cost 参数和每核心权重数组提供，但计算每个核心成本的算术(LatencyEstimator / AssignQueueIDsToAsyncStart / AssignQueueIDsForComputation 中的队列占用提要)未被追踪 - 抢七在结构上是固定的(成本上升)，但不是封闭形式。
预留预算初始值设定项。 线程本地每个资源预算递减并门控 >= 2，但播种 TLS long 及其初始值(NumEmbeddingDevices 预留 vs 每个资源固定核心数上限 2 vs LogicalDevicesPerChip(SC))的编写器是没有追踪到初始化站点。
资源类型 → 物理核心注册表。 GetSparseCoreResources 返回 {0, 23..28} 资源类型 ID 和 GetAllowedCores 将芯片分组到其下，但未枚举将资源类型 ID 映射到它可以占用的一组物理 SC 核心的全局表(以及超级核心配对如何将两个核心折叠到一个资源插槽中)。

组件	关系
`SparseCoreCompiler::RunHloScheduler` (`0x1306F820`)	拥有卸载门；当它成立时运行 SC 延迟隐藏调度程序
`SparseCoreQueueAssignment`	拥有 `GetAllowedCores` + `SelectCores` 的通证(本页政策)
`AssignQueueIDsToAsyncStart` (`0x10FDF480`)	每个集合驱动程序：调用 `GetAllowedCores` → `SelectCores` → `__sort` → 原型写入
`GetSparseCoreConfig` (`0x1C868D20`)	读取控制路由和资源选择的 `SparseCoreConfig.offload` 枚举

交叉引用

On-Pod 集体 — 剖面图 — 衬底分裂和门在选择 SC 调度器中的作用(§5)；此页面是门位+选择策略详细信息
SC内核选择 — 字节级 GetAllowedCores / SelectCores 主体(Info 结构、SIMD 搜索、HloReachabilityMap 探针、瑞士表掩码构建)本页总结
物理核心放置 — 下游数字 __sort 以及映射到 physical_core_indices 的每颜色→物理核心索引
SC-卸载配置生成器 — ConstructConfigForCollectiveUniDirNDGroups<*> 和 *OffloadConfig 原型是门控集体生产的
分层种类 — 卸载集体配置生成器内的扁平与分层阶段划分
张量分割/ND 平面 — tensor_split_factor / ND 平面推导，消耗所选内核的 count
SparseCoreConfig(GetSparseCoreConfig) — 完整的 SparseCoreConfig 原型字段映射和 Offload 枚举描述符字节
SC队列分配和预留 — [this+0xC0] 预留 btree_map(实时线程本地预算 §3 强制执行的残留孪生)
资源类型分类 — 这些资源索引到的 {0, 23..28} 调度资源类型空间
返回索引

SC 内核选择和卸载门 ​

摘要 ​

1. 卸载门 ​

用途 ​

算法 ​

两个目标位 ​

每代硬件基础 ​

2. SparseCoreConfig.offload — op 类型分类器 ​

用途 ​

枚举所在的位置 ​

3. GetAllowedCores — 候选核心掩码 ​

用途 ​

两个输入 ​

五个谓词 ​

抢七局 ​

5. 验证说明 ​

我们还没有什么 ​

相关组件 ​

交叉引用 ​

SC 内核选择和卸载门

摘要

1. 卸载门

用途

算法

两个目标位

每代硬件基础

2. `SparseCoreConfig.offload` — op 类型分类器

用途

枚举所在的位置

3. `GetAllowedCores` — 候选核心掩码

用途

两个输入

五个谓词

抢七局

5. 验证说明

我们还没有什么

相关组件

交叉引用