SC 队列分配与预留

本页上的每个地址、偏移、字段号和常量都从 libtpu-0.0.40-cp314 wheel 中的 libtpu.so 逐字节读取（BuildID md5 89edbbe81c5b328a958fe628a9f2207d，构建 libtpu_lts_20260413_b_RC00）。.text VMA == 文件偏移（0xE63C000）；.rodata VMA == 文件偏移（0x84A0000）；.data.rel.ro VMA - 0x200000 == 文件偏移。其他版本会不同。

摘要

本页是 xla::jellyfish::SparseCoreQueueAssignment 在成员偏移 [this+0xC0] 暂存的按资源预留映射：一个以九个 SparseCore / 异步集合调度资源类型 ID {2, 3, 6, 23, 24, 25, 26, 27, 28} 为键的 absl::btree_map<long resource_id, long limit>，每个键携带一个来自具名 TpuCompilationEnvironment (TCE) 编译开关的并发 / 重叠限制。该映射是 SparseCore 后端对“每类资源允许同时在途多少个”的意图表述。本页记录三件事，全部以字节为锚点：

预留 btree_map<long,long> — 它的九个键、每个键的值来源（pass 对象的 +0x18..+0x58 窗口），以及它在哪里构建（RunImpl）和释放（析构函数）。
为每个预留值播种的 TCE AutoOr<long> 字段标签 — 字段号 {288, 924, 925, 1088, 1089, 1090, 1091, 1092}（键 28 是硬编码 INT64_MAX）、它们的 proto 字段名，以及构建点处的 AUTO -> 1 回退极性。
读取 [this+0xC0] 的真实消费者 — 一个明确的否定结果。SparseCoreQueueAssignment 的任何成员都不读取该映射。它只被构建一次并销毁；分配路径（GetAllowedCores / SelectCores / AssignQueueIDsToAsyncStart）从不查询它。该映射在 v0.0.40 中是一个已暂存 / 未接线成员：它保存配置的按资源上限，但不门控核心选择。相同上限的执行位于别处 — LatencyHidingScheduler 的 AsyncTracker 资源限制表中，该表在相同资源 ID 下从相同 TCE 字段构建（见 ResourceType 分类）。

重新实现者需要的心智模型是，这个 [this+0xC0] 映射是实时调度器并发表的残留副本，不是缺失功能。预留值会从具名开关完整且正确地解析；只有读取者在 SparseCoreQueueAssignment 翻译单元中不存在。

对于重新实现，契约是：

预留映射是位于 [this+0xC0] 的 btree_map<long,long>，使用 map_params_impl<long,long> 节点参数和共享空节点哨兵 0x2177A4B0。键是资源类型 ID；值是并发 / 重叠限制。
九个固定键，按构建顺序：{3, 6, 23, 24, 25, 26, 27, 28} 加上键 2。 每个值都从 pass 对象的 +0x18..+0x58 qword 窗口读取，该窗口由 AddPass<...> 构造函数从调用方的 SparseCoreQueueAssignmentConfig 栈结构体用 vmovups 复制而来。
九个值中有七个来自 AutoOr<long> TCE 开关（此处 AUTO -> 1）。 键 2 是原始 proto3 int32（默认 0）；键 28 是硬编码的 INT64_MAX 字面量（无 TCE 字段）。七个 AutoOr<long> 读取在此构建点使用 AUTO 未设置 -> 1 的极性 — 不同于实时调度器的 AUTO -> INT64_MAX。
该映射在 v0.0.40 中没有消费者。 它在 RunImpl（0x10FE4000）中构建一次，并在 ~SparseCoreQueueAssignment（0x10FE4BA0）中释放。pass 区间内任何位置都不存在对 [this+0xC0] 的 internal_find / lower_bound / operator[]。预留由此 pass 暂存，而不是执行。
执行路径是 AsyncTracker 并发表，不是这个映射。 相同键 {2,3,6,23..28} 和相同字段 {288,924,925,1088..1092} 会再次由 GetTpuAsyncTracker / GetSchedulerConfig 读取，并由 TpuAsyncTracker::GetNumAvailableResources（0x10FFF600）转为共同发射上限。


预留映射	`[SparseCoreQueueAssignment + 0xC0]` 处的 `btree_map<long,long>`
节点参数	`map_params_impl<long,long>`；空节点哨兵 `0x2177A4B0`
键（资源 ID）	`{2, 3, 6, 23, 24, 25, 26, 27, 28}`
值窗口	pass 对象 `+0x18..+0x58`（9 个 qword；从 `Config[+0x00..+0x40]` 复制）
构建位置	`SparseCoreQueueAssignment::RunImpl`（`0x10FE4000`）— 九次 `insert_hint_unique`
插入器	`btree<map_params_impl<long,long>>::insert_hint_unique`（`0x10FE6E40`）
根存储	`*(this+0xC0) = built-tree-root`（`@0x10FE42B9`）
释放位置	`~SparseCoreQueueAssignment::D2`（`0x10FE4BA0`）-> `btree_node::clear_and_delete`（`0xF7D0400`）
配置构建器	`(anon)::RunHloScheduler`（`0x109718C0` 区域）— 读取 TCE，填充 `Config[-0x138]`
配置 -> pass 复制	`HloPassPipeline::AddPass<SparseCoreQueueAssignment,...>`（`0x10975FC0`）— 三次 `vmovups`
映射消费者	v0.0.40 中无（明确否定；见下方单元）
实时执行者（别处）	`TpuAsyncTracker::GetNumAvailableResources`（`0x10FFF600`）— 见 ResourceType 分类

`[this+0xC0]` 处的预留 `btree_map<long,long>`

目的

SparseCoreQueueAssignment 为每种资源类型预留一个槽位：存储在 pass 对象 +0xC0 的 btree_map<long resource_id, long limit>。键是集合可占用的 AsyncTracker / SparseCore 调度资源类型 ID（与 GetSparseCoreResources 产生的 {0,23..28} 空间相同，加上基础异步集合 ID {2,3,6}）；值是编译环境配置的按资源并发或重叠限制。意图是对每类资源可排队数量设置按资源上限；v0.0.40 中的实现是一个构建后释放、没有读取者的成员（见下方映射消费者）。

映射结构

该映射是一个带 map_params_impl<long,long> 的 absl::btree_map — 反编译中构建和释放点引用的 btree<...map_params_impl<long,long>...> 模板实例逐字节确认了这一点。新建 / 空映射的根指向共享空节点哨兵 EmptyNode()::empty_node（重定位地址 0x2177A4B0）；构造函数用该哨兵初始化 [this+0xC0]，并将 [this+0xC8]（大小 / 最右叶字段）置零。

text

SparseCoreQueueAssignment object (sizeof 0xF8, _Znwm 0xF8 @0x10975FD5):
  [this+0x00]  vptr (= 0x2181D8C8 + 0x10, @0x10976020)
  [this+0x08]  Target*            (= r12 @0x10976023)
  [this+0x10]  unique_ptr<LatencyEstimator>
  [this+0x18 .. +0x58]  the 9-qword reservation VALUE window  (Config copy; see below)
  [this+0x60]  bool   (the 10th Config field; NOT a map value — see callout)
  ...
  [this+0xC0]  btree_map<long,long> root pointer  ← the RESERVATION MAP
  [this+0xC8 .. +0xD7]  btree size / rightmost-leaf fields
```text

### 算法 — 构建（`RunImpl`）

`RunImpl`（`0x10FE4000`）组装九个 `{key, value}` 栈对，并通过 `insert_hint_unique`（`0x10FE6E40`）逐个插入一个新的 btree，然后将构建好的树根移动到 `[this+0xC0]`。反编译逐字显示了这些 pair 的设置（`a2` = pass 的 `this`）：

```c
// SparseCoreQueueAssignment::RunImpl  @0x10FE4000  (a2 = this)
v80    = *(_QWORD *)(a2 + 24);   //  key 2  value  ← obj[+0x18]   (Config[+0x00])
v81[0] = 3;  v81[1] = *(_QWORD *)(a2 + 32);   //  key 3  ← obj[+0x20]  (Config[+0x08])
v82[0] = 6;  v82[1] = *(_QWORD *)(a2 + 40);   //  key 6  ← obj[+0x28]  (Config[+0x10])
v83[0] = 23; v83[1] = *(_QWORD *)(a2 + 48);   //  key 23 ← obj[+0x30]  (Config[+0x18])
v84[0] = 24; v84[1] = *(_QWORD *)(a2 + 56);   //  key 24 ← obj[+0x38]  (Config[+0x20])
v85[0] = 25; v85[1] = *(_QWORD *)(a2 + 64);   //  key 25 ← obj[+0x40]  (Config[+0x28])
v86[0] = 26; v86[1] = *(_QWORD *)(a2 + 72);   //  key 26 ← obj[+0x48]  (Config[+0x30])
v87[0] = 27; v87[1] = *(_QWORD *)(a2 + 80);   //  key 27 ← obj[+0x50]  (Config[+0x38])
v88[0] = 28; v88[1] = *(_QWORD *)(a2 + 88);   //  key 28 ← obj[+0x58]  (Config[+0x40])

v100 = &btree<map_params_impl<long,long>>::EmptyNode()::empty_node;   // fresh root
// nine insert_hint_unique<long, pair<long,long> const&>(&tree, &root, ...):
btree<…>::insert_hint_unique(&v96, &v100, v101, *(u8*)(v101+10), v80_pair, …);   // key 2
btree<…>::insert_hint_unique(&v96, &v100, v101, *(u8*)(v101+10), v81, v81);      // key 3
// … keys 6,23,24,25,26,27,28 …
*(_QWORD *)(a2 + 192) = v100;    // root → [this+0xC0]   (@0x10FE42B9)

算术逐字节对齐：a2 + 24 = obj[+0x18]，a2 + 88 = obj[+0x58]，而 a2 + 192 = [this+0xC0]。这九个值就是 AddPass 构造函数从调用方 SparseCoreQueueAssignmentConfig 复制来的 qword 窗口（见下一单元）。

注意 — 键 2 构建为裸 pair v80（第一次插入），而键 {3,6,23..28} 是 vNN[0]=key; vNN[1]=value pair。九个值都通过同一个 insert_hint_unique（0x10FE6E40）流入同一个根。构建是无条件的 — 没有分支会跳过某个键 — 因此 RunImpl 之后该映射始终正好有九个条目。

算法 — 释放（析构函数）

~SparseCoreQueueAssignment::D2（0x10FE4BA0）通过在 [this+0xC0] 上调用 btree_node::clear_and_delete（0xF7D0400）释放该映射：

// ~SparseCoreQueueAssignment::D2  @0x10FE4BA0
btree_node<map_params_impl<long,long>>::clear_and_delete(*((_QWORD **)this + 24));  // [this+0xC0]
// ( *((_QWORD*)this + 24) == this+0x180? no — index 24 * 8 = 0xC0 )
```text

`*((_QWORD*)this + 24)` 是 `this + 24*8 = this + 0xC0`。析构函数还会释放该 pass 在 `[this+0xA0]` 和 `[this+0x80]` 持有的两个 Swiss table（来自 `GetAllowedCores` 的 `flat_hash_set`/分组状态），但只有 `[this+0xC0]` 的 `clear_and_delete` 会触碰预留映射。

### 函数映射

| 函数 | 地址 | 作用 |
|---|---|---|
| `SparseCoreQueueAssignment::RunImpl` | `0x10FE4000` | 构建九条目预留映射；根 -> `[this+0xC0]` |
| `btree<map_params_impl<long,long>>::insert_hint_unique` | `0x10FE6E40` | 按键插入器（x9） |
| `btree<map_params_impl<long,long>>::EmptyNode` | `0x2177A4B0` | 新根的空节点哨兵 |
| `~SparseCoreQueueAssignment::D2` | `0x10FE4BA0` | 通过 `clear_and_delete` 释放映射 |
| `btree_node<map_params_impl<long,long>>::clear_and_delete` | `0xF7D0400` | 递归 btree 释放 |

---

## 为值播种的 TCE 字段标签

### 目的

每个预留值都源自具名 `TpuCompilationEnvironment` (TCE) proto 字段。数据路径有三跳，全部逐字节确认：调用方读取 TCE 并在栈上构建 `SparseCoreQueueAssignmentConfig`；`AddPass` 构造函数用 `vmovups` 将该 Config 复制进 pass 对象的 `+0x18..` 窗口；`RunImpl` 将该窗口作为九个映射值读取。本单元命名每个字段 — 它的 proto **字段号**（从 `TpuCompilationEnvironment::_InternalSerialize`，`0x1DB41DC0` 解码）、它的**字段名**（来自切出的 `FileDescriptorProto`）、它的**类型**，以及它在此构建点的 **AUTO 回退**。

### 跳 A — Config 构建（调用方）

调用方（`RunHloScheduler` 内部的 SparseCore 调度器设置，区域 `0x109718C0..0x10971A99`）通过 `GetTpuCompEnv`（`0x1D73DD20`，`GetMutableEnv<TpuCompilationEnvironment>` `_impl_`）取得 TCE，并在 `[rbp-0x138]` 组装 `SparseCoreQueueAssignmentConfig`：

```text
key 2  (int32, raw):   movslq 0xF78(TCE),rax        → Config[+0x00]    field 288
keys 3,6,23..27 (AutoOr<long>):
   rdi = TCE[+OFF]; if null → AutoProto_globals_ (0x223C8968, cmove %r15);
   AutoOr<long>::FromProtoOrDie (0x1092F7E0) → {value rax, has-bit dl};
   test $1,%dl ; mov $1,%r13d ; cmove %r13,%rax    ⇒ AUTO (oneof unset) → 1
       TCE+0x460 → Config[+0x08]   field 924
       TCE+0x468 → Config[+0x10]   field 925
       TCE+0x940 → Config[+0x18]   field 1088
       TCE+0x948 → Config[+0x20]   field 1089
       TCE+0x950 → Config[+0x28]   field 1090
       TCE+0x958 → Config[+0x30]   field 1091
       TCE+0x960 → Config[+0x38]   field 1092
key 28 (constant):    movabs $0x7FFFFFFFFFFFFFFF   → Config[+0x40]    (no TCE field — INT64_MAX)

跳 B — `AddPass` 的 `vmovups` 复制

HloPassPipeline::AddPass<SparseCoreQueueAssignment,...>（0x10975FC0）分配 pass（_Znwm 0xF8），并用三次重叠的 32 字节移动将 Config 参数（rcx/r14）复制到对象：

text

vmovups (%r14),%ymm0      → obj[+0x18..+0x37]   (= Config[+0x00..+0x1F])
vmovups 0x20(%r14),%ymm1  → obj[+0x38..+0x57]   (= Config[+0x20..+0x3F])
vmovups 0x30(%r14),%ymm2  → obj[+0x48..+0x67]   (= Config[+0x30..+0x4F])
```text

因此 `obj[+0x18] = Config[+0x00]`，……，`obj[+0x58] = Config[+0x40]`（重叠的 `obj[+0x48..+0x57]` 由 `ymm1` 和 `ymm2` 写入相同内容）。构造函数还会存储 `Target` 指针（`obj[+0x08]`）、`LatencyEstimator` 的 `unique_ptr`（`obj[+0x10]`）、vtable（`0x2181D8C8+0x10`），并设置空映射根（`obj[+0xC0] = empty_node 0x2177A4B0`）。

### 完整字段表

`TCE _impl_ off` 列是来自 `RunHloScheduler` 的字节偏移；`TCE field name` 列是匹配的 `AbslFlagDefaultGenFor<name>` 符号中的名称。`TCE field #` 列是 proto 源标签（见表后说明）。

| 映射键 | obj 偏移 | Config 偏移 | TCE `_impl_` 偏移 | TCE 字段号 | TCE 字段名 | Proto 类型 | AUTO 回退（此处） |
|---|---|---|---|---|---|---|---|
| 2 | `+0x18` | `+0x00` | `0xF78` | 288 | `xla_max_concurrent_async_all_gathers` | int32 (raw) | proto3-zero `0` |
| 3 | `+0x20` | `+0x08` | `0x460` | 924 | `xla_max_concurrent_async_all_reduces` | `AutoOr<long>` | `1` (cmove) |
| 6 | `+0x28` | `+0x10` | `0x468` | 925 | `xla_max_concurrent_async_reduce_scatters` | `AutoOr<long>` | `1` (cmove) |
| 23 | `+0x30` | `+0x18` | `0x940` | 1088 | `xla_tpu_sparse_core_gather_overlap_limit` | `AutoOr<long>` | `1` (cmove) |
| 24 | `+0x38` | `+0x20` | `0x948` | 1089 | `xla_tpu_sparse_core_scatter_overlap_limit` | `AutoOr<long>` | `1` (cmove) |
| 25 | `+0x40` | `+0x28` | `0x950` | 1090 | `xla_tpu_sparse_core_data_formatting_overlap_limit` | `AutoOr<long>` | `1` (cmove) |
| 26 | `+0x48` | `+0x30` | `0x958` | 1091 | `xla_tpu_sparse_core_kernel_overlap_limit` | `AutoOr<long>` | `1` (cmove) |
| 27 | `+0x50` | `+0x38` | `0x960` | 1092 | `xla_tpu_sparse_core_sort_overlap_limit` | `AutoOr<long>` | `1` (cmove) |
| 28 | `+0x58` | `+0x40` | — | — | （硬编码 `INT64_MAX`；不是 TCE 字段） | 常量 | `0x7FFFFFFFFFFFFFFF` |

每个字段的 `_impl_` 字节偏移都由 `RunHloScheduler` Config 构建区域逐字节确认（`GetTpuCompEnv(...) + N` 加载：`+3960`、`+1120`、`+1128`、`+2368`、`+2376`、`+2384`、`+2392`、`+2400` — 见跳 A），每个偏移都绑定到一个不同的 `AutoOr<long>` proto 字段，其 flag 符号名（`AbslFlagDefaultGenFor<name>`）匹配下表。proto **字段号**记录自 `TpuCompilationEnvironment::_InternalSerialize`（`0x1DB41DC0`）；其序列化器 tag 多数是预移位 varint，而不是裸 `mov $imm` 立即数，因此偏移↔名称绑定是重新实现者需要的内容，编号是 proto 源标签。`925` tag（`0x39D`）可在序列化器主体中观察到；int32 键 2 携带 wire tag `0x1280`（varint `0x1280` -> 2304 -> 字段 288，wire-type 0）。

> **注意 — AUTO 极性是调用点局部的。** 此处七个 `AutoOr<long>` 读取会将*未设置*（AUTO）oneof 解析为 **`1`**（`test $1,%dl ; cmove $1` 形态）。这只是 `AddPass` Config 构建点的极性。*实时*调度器在 `GetTpuAsyncTracker`（`0x10975520`）中读取**相同**五个 SC 重叠开关（`1088..1092`）时，AUTO -> **`INT64_MAX`**（无上限）— 相同字段，相反默认值，不同调用点。重新实现者不能假定这些开关有一个全局默认值。执行路径中的语义默认值是“无上限”；此处的 `1` 没有实际效果，因为该映射没有读取者。
>
> **陷阱 — 第 10 个 Config 字段不是映射值。** Config 还携带第十个字段，即来自 `AutoOr<bool>(TCE[+0xC88])` 的 bool（TCE 字段 **1202** `xla_tpu_rerun_latency_hiding_scheduler_post_sc_assignment`，`and $0x101; cmp $0x101; setne` AUTO-on 惯用法）。它被复制到 `obj[+0x60]`，位于 `+0x18..+0x58` 映射值窗口*之外*，`RunImpl` 不会为映射读取它。它门控的是 SC 分配后的 `LatencyHidingScheduler` 重新运行，而不是任何预留值。见下方*实时执行者*。

### 资源 ID -> 集合对应关系

这些键是 AsyncTracker 调度资源类型 ID。基础 ID `{2,3,6}` 和 SC ID `{23..28}` 索引的资源空间与 `GetSparseCoreResources`（`0x10FDC0A0`）和 `MayAddSparseCoreResource`（`0x11000480`）产生的相同。权威 `kSparseCore*` 名称和 opcode->id switch（包括 `0x56 -> 12`、`0x5d -> 6` 映射）位于 [ResourceType 分类](../sched/scheduler-resourcetype-model.md)；开关名称所归属的 SC 类别为：

| 键 | 资源类别 | TCE 开关（字段） | 单位 | 构建处默认值 |
|---|---|---|---|---|
| 2 | 基础异步集合 | `xla_max_concurrent_async_all_gathers` (288) | max-concurrent | 0 |
| 3 | 基础异步集合 | `xla_max_concurrent_async_all_reduces` (924) | max-concurrent | 1 (AUTO) |
| 6 | 基础异步集合 | `xla_max_concurrent_async_reduce_scatters` (925) | max-concurrent | 1 (AUTO) |
| 23 | `kSparseCoreGather` | `xla_tpu_sparse_core_gather_overlap_limit` (1088) | overlap-limit | 1 (AUTO) |
| 24 | `kSparseCoreScatter` | `xla_tpu_sparse_core_scatter_overlap_limit` (1089) | overlap-limit | 1 (AUTO) |
| 25 | `kSparseCoreDataFormatting` | `xla_tpu_sparse_core_data_formatting_overlap_limit` (1090) | overlap-limit | 1 (AUTO) |
| 26 | `kSparseCoreKernel` | `xla_tpu_sparse_core_kernel_overlap_limit` (1091) | overlap-limit | 1 (AUTO) |
| 27 | `kSparseCoreSort` | `xla_tpu_sparse_core_sort_overlap_limit` (1092) | overlap-limit | 1 (AUTO) |
| 28 | 兜底 | （无 — 硬编码 `INT64_MAX`） | unlimited | `INT64_MAX` |

> **注意 — 名称与产生 opcode。** 值来源*名称*（`all_gathers` -> 键 2，`all_reduces` -> 键 3，`reduce_scatters` -> 键 6）并不与每个资源类型的产生 opcode 名称 1:1 匹配；绑定依据是 Config 作者选择的数值资源类型 ID，也就是 `RunImpl` 编码的内容。名称记录的是逐字节精确的值来源，而不是重新归因给 opcode。opcode->resource-type switch 位于 [ResourceType 分类](../sched/scheduler-resourcetype-model.md)。

### 函数映射

| 函数 | 地址 | 作用 |
|---|---|---|
| `(anon)::RunHloScheduler`（Config 构建区域） | `0x109718C0..0x10971A99` | 读取 TCE；构建 `SparseCoreQueueAssignmentConfig[-0x138]` |
| `GetTpuCompEnv` | `0x1D73DD20` | 返回 `TpuCompilationEnvironment` `_impl_` |
| `AutoOr<long>::FromProtoOrDie` | `0x1092F7E0` | 打包 `{value rax, has-bit dl}` 读取器 |
| `HloPassPipeline::AddPass<SparseCoreQueueAssignment,...>` | `0x10975FC0` | 三次 `vmovups`；Config -> `obj[+0x18..]` |
| `TpuCompilationEnvironment::_InternalSerialize` | `0x1DB41DC0` | 字段号解码来源 |
| `AutoProto_globals_` | `0x223C8968` | null-TCE 回退默认实例 |

---

## 映射消费者 — 明确否定结果

### 发现

**`SparseCoreQueueAssignment` 的任何成员都不读取 `[this+0xC0]`。** 预留映射只构建一次（`RunImpl`）并释放一次（析构函数）；两者之间没有任何内容查询它。对 pass 区间（`0x10FDA3C0..0x10FE4BA0`）逐成员完整扫描 `internal_find` / `lower_bound` / `operator[]` 对 `[this+0xC0]` 的使用，以及任何会将映射按引用传给辅助函数的 `lea 0xC0(this)` / `add $0xC0, this`，只发现了 `RunImpl` 构建和析构函数释放。该预留在 v0.0.40 中是一个**已暂存 / 未接线**成员：它携带配置的按资源上限，但在发布的核心选择路径中不排除、不限制、也不划分候选核心。

### 反编译交叉检查

四个选择路径成员加上按计算驱动函数都已扫描；没有一个读取 `[this+0xC0]`：

| 成员函数 | 地址 | 行数 | `[this+0xC0]` 读取 |
|---|---|---|---|
| `GetAllowedCores` | `0x10FDA3C0` | 1511 | 对 `this` 为 **0**（反编译第 748 行的一个 `+192` 位于本地设备遍历指针 `[hlo+0xC0]`，不是 `this`） |
| `SelectCores` | `0x10FDC4E0` | — | 0 |
| `AssignQueueIDsToAsyncStart` | `0x10FDF480` | 833 | **0** |
| `GetLogicalAssignmentGroups` | `0x10FE0820` | — | 0 |
| `AssignQueueIDsForComputation` | `0x10FE1D20` | — | 0 |
| `RunImpl` | `0x10FE4000` | 557 | 仅构建（根存储 `@0x10FE42B9`） |
| `~SparseCoreQueueAssignment::D2` | `0x10FE4BA0` | — | 仅释放（`clear_and_delete`） |

> **陷阱 — `[hlo+0xC0]` 不是 `[this+0xC0]`。** `GetAllowedCores` 确实携带一个 `+0xC0` 访问（反编译第 748 行：`v199 = (long*)((char*)v220 + 192)`），但 `v220` 是一个指向设备分配 / replica-group 遍历的*本地*指针 — 是对指令结构的第二次遍历 — **不是** pass 的 `this`。那里的反汇编 `add $0xC0`（`@0x10FDAE39`）操作的是 `[hlo+0xC0]`，即设备分配遍历。任何位置都没有 `[this+0xC0]`（预留映射）读取。这两个 `0xC0` 偏移的碰撞只是布局巧合，不是该映射的使用。

### 实际排除候选核心的内容

在 `GetAllowedCores` 中真正运行的候选核心排除是 Swiss-table 分组（位于 `0x2181D940` 的 `flat_hash_map<resource_id, btree_set<chip_id>>` 和位于 `0x21639C10` 的按芯片占用 `flat_hash_map<chip_id, long>`）加上按资源线程局部预算门控（`__tls_get_addr(&qword_22048D78)`，门控 `>= 2`）。这个机制 — 不是此映射 — 是选择路径中活跃的预留式过滤器；见 [SC 核心选择](sc-core-selection.md)。那里的线程局部预算被记录为独立的、逐字节确认的门控；此 `[this+0xC0]` 映射*不是*它的后备存储（预算播种点在该页仍未解决，但它不是从此映射播种的）。

### 实时执行者 — 相同上限在哪里生效

预留映射是死的，但它暂存的按资源上限在**别处执行** — 位于 `LatencyHidingScheduler` 的 `AsyncTracker` 资源限制表中，该表在**相同**资源 ID 下从**相同** TCE 字段构建：

- **基础集合 `{2,3,6}` <- 字段 `{288,924,925}`**，经由 `GetSchedulerConfig` -> `SchedulerConfig[+0x20/+0x28/+0x30]` -> `TpuAsyncTracker` ctor -> `tracker[+0x68/+0x70/+0x78]` -> `AsyncTracker::SetConcurrentResourceLimits`（`0x13615800`）键 `{2,3,6}`。
- **SC 类别 `{23..27}` <- 相同偏移 `+2368..+2400`（字段 `{1088..1092}`）** 在 `GetTpuAsyncTracker`（`0x10975520`，第 118-151 行）中再次读取 — 但此处 AUTO -> `INT64_MAX`，逐字节确认 — 再加上 TCE 偏移 `+2696`（`0xA88`）处的一个额外 `AutoOr<long>` 开关；这些值流经 `TpuAsyncTracker::Create`（第 179 行）进入 tracker，并由 `TpuAsyncTracker::GetNumAvailableResources`（`0x10FFF600`）按资源 ID 输出。

调度器会拒绝共同发射超过 `limit` 个给定资源类型的异步 op。字段 1202 门控（第 10 个 Config bool）控制 SparseCore 队列分配之后的*第二次* LHS pass，它会构建一个新的 `TpuAsyncTracker` 并读取**相同**重叠限制 — 因此这些上限在队列分配前后都会应用。本页只记录死的 `SparseCoreQueueAssignment` 映射；实时表位于 [ResourceType 分类](../sched/scheduler-resourcetype-model.md) 和 [LatencyHidingScheduler 核心](../sched/latency-hiding-scheduler-core.md)。

> **注意 — 是暂存副本，不是缺失功能。** `[this+0xC0]` 映射使用的*相同* `{2,3,6,23..28}` 来自实时 `AsyncTracker` 表使用的*相同*字段 `{288,924,925,1088..1092}`。它是调度器并发表的冗余 / 残留副本，解析正确，但此 pass 中没有任何人读取。重新实现者可以为保真构建该映射，但应理解执行发生在调度器中，并且在*执行*路径中未设置 SC 重叠开关的运行时默认值是 `INT64_MAX`（无上限），不是此构建点暂存的 `1`。

### 函数映射

| 函数 | 地址 | 作用 |
|---|---|---|
| `SparseCoreQueueAssignment::GetAllowedCores` | `0x10FDA3C0` | 候选掩码构建；读取 `[hlo+0xC0]`，不是 `[this+0xC0]` |
| `SparseCoreQueueAssignment::AssignQueueIDsToAsyncStart` | `0x10FDF480` | 按集合驱动；无映射读取 |
| `TpuAsyncTracker::GetNumAvailableResources` | `0x10FFF600` | 实时按资源共同发射上限（真实执行者） |
| `AsyncTracker::SetConcurrentResourceLimits` | `0x13615800` | 构建实时 `resource_type->limit` 映射 |
| `(anon)::GetTpuAsyncTracker` | `0x10975520` | 为 tracker 读取 `1088..1092`/`1130`（AUTO -> `INT64_MAX`） |

---

## 分代说明

预留映射的构建、字段读取或（不存在的）消费者在代码中都没有按代分支。该映射始终有九个条目；构建点的 AUTO -> 1 极性是常量；资源类型 ID 是常量。*执行*调度器路径会因某些 SC-offload 并发默认值按 `TpuVersion` 分支（`ShouldEnableConcurrentSparseCoreOffloading` 基础，`TpuVersion == 5`），但那是实时 `AsyncTracker` 路径，不是此映射；见 [GetSparseCoreConfig](getsparsecoreconfig.md) 获取按代 SC-offload 门控。该映射的存在和它的九个具名开关记录了跨硅片的*预期*预留表面；没有读取者则记录了此构建中执行并未通过该 pass 接线。

---

## 相关组件

| 名称 | 关系 |
|---|---|
| `SparseCoreQueueAssignment::RunImpl`（`0x10FE4000`） | 构建预留映射；根 -> `[this+0xC0]` |
| `HloPassPipeline::AddPass<SparseCoreQueueAssignment,...>`（`0x10975FC0`） | 将 Config 结构复制到 pass 对象的值窗口 |
| `(anon)::RunHloScheduler`（`0x109718C0` 区域） | 读取 TCE 开关；构建 Config 结构 |
| `GetSparseCoreResources`（`0x10FDC0A0`） | 产生该映射所键入的相同 `{0,23..28}` 资源类型空间 |
| `TpuAsyncTracker::GetNumAvailableResources`（`0x10FFF600`） | 相同上限的实时消费者（真实执行者） |
| `AsyncTracker::SetConcurrentResourceLimits`（`0x13615800`） | 实时 `resource_type->limit` 表构建器 |

## 交叉引用

- [SC 核心选择](sc-core-selection.md) — 此映射本应供给的 `GetAllowedCores` / `SelectCores` 策略；`AssignQueueIDsToAsyncStart` 调用链，以及实际会排除核心的线程局部预算（不是此映射）。
- [GetSparseCoreConfig](getsparsecoreconfig.md) — offload op-type 枚举和 SC-offload 调度器门控；`{23..28}` 资源类型类别的来源。
- [SC 后端流水线](sc-backend-pipeline.md) — 队列分配 pass 所在的 SC-MLO pass 流水线。
- [SparseCore 硬件架构](architecture.md) — 几何结构和 4:1 SC:TC 比例，它限定策略可选择的物理核心数。
- [SparseCore 概览](overview.md) — 第 IX 部分的导航入口。
- [ResourceType 分类](../sched/scheduler-resourcetype-model.md) — **实时** `AsyncTracker` 资源限制表：`SetConcurrentResourceLimits`、`GetNumAvailableResources`、`kSparseCore*` 名称，以及实际执行的相同 `{2,3,6,23..28}` 键 / `{288,924,925,1088..1092}` 字段。
- [LatencyHidingScheduler 核心](../sched/latency-hiding-scheduler-core.md) — 其 `AsyncTracker` 实际限制并发异步发射的调度器，以及字段 1202 门控的 SC 分配后重新运行。
- **二进制：** `extracted/libtpu-0.0.40-cp314-cp314-manylinux_2_31_x86_64/libtpu/libtpu.so`（build-id `89edbbe81c5b328a958fe628a9f2207d`）
- **索引项：** 第 IX 部分 — SparseCore & BarnaCore / SparseCore 后端 — [返回索引](../index.md)

SC 队列分配与预留 ​

摘要 ​

[this+0xC0] 处的预留 btree_map<long,long> ​

目的 ​

映射结构 ​

算法 — 释放（析构函数） ​

跳 B — AddPass 的 vmovups 复制 ​

SC 队列分配与预留

摘要

`[this+0xC0]` 处的预留 `btree_map<long,long>`

目的

映射结构

算法 — 释放（析构函数）

跳 B — `AddPass` 的 `vmovups` 复制