ResourceType 分类体系

地址适用于来自 libtpu-0.0.40-cp314 wheel 的 libtpu.so（build-id 89edbbe81c5b328a958fe628a9f2207d）。其他版本会有所不同。所有 .text/.rodata 地址都是虚拟地址；对该二进制而言，.text VMA == 文件偏移 0xe63c000，.rodata VMA == 文件偏移 0x84a0000。

摘要

LatencyHidingScheduler 通过询问每个异步 op 消耗哪种物理资源以及硬件拥有多少该资源，来决定两个异步 op 是否可以共同发射。资源名形成一个单一扁平整数枚举 ResourceType，范围为 0..46（47 个 ID）。前 13 个是标准 XLA AsyncTracker collective 类（kNoResource、kAllReduce、kReduceScatter 等）；其余 34 个是 jellyfish::TpuAsyncTracker 添加的 TPU 目标资源（六个有方向的 ICI ring link、DCN 带宽、两个 host-DMA tap、六个 SparseCore 引擎类、VMEM，以及十六个用户 custom-collective lane）。本页从两个 GetResourceName 查找路径、GetResourceTypeForOp opcode→id switch 和对 HLO op 分类的六个 MayAdd* 生产者、GetNumAvailableResources 的按资源并发上限来源表、GetResourceHazardType overlap-class 表，以及决定哪个 AsyncTracker 子类为 computation 建模的按 pass tracker 选择中，恢复完整的 enum value→name 表。

读者不能把这个 ResourceType enum 与 Resource Enum 记录的 cost-model ResourceVector::Resource enum 混淆。它们是同一二进制中两个不同的“resource”抽象。cost-model Resource（23 个槽位，嵌套于 ResourceVector）度量cycle weight，也就是功能单元繁忙多少 cycle，并馈送 MaxResourceCycles。此 ResourceType（47 个 ID，一个 AsyncTracker enum）度量并发限制，也就是某类异步 op 可同时在途多少个，并馈送 GetNumAvailableResources。它们只在一个物理概念上重叠（六个 ICI ring link：cost-model 槽位 R[13..18] 对 scheduler ids 14..19），但它们是独立 enum，拥有独立 value space、独立 name table 和独立 consumer。本页记录 ResourceType；cost-model Resource 是它的兄弟。

对于重新实现，契约是：

• 47-ID ResourceType enum：来自 AsyncTracker::GetResourceName 的 base XLA {0..12}，以及来自 TpuAsyncTracker::GetResourceName 的 jellyfish target {13..46}，其中有两个未命名 catch-all（ids 28 和 46），ids 30..45 共享 kCustomCollective 字符串。
• GetResourceTypeForOp，即 base opcode→id switch（规范 XLA collective map），加上 jellyfish 编排器 GetResourcesFromInstructionImpl 及其六个在其上添加 target ids 的 MayAdd* 生产者。
• GetNumAvailableResources，即按 id 的并发上限及其来源：TCE knob、一个硬件 core-count（id 22，CoresPerChip(SC)/LDPC(SC)）和常量。
• GetResourceHazardType，即每个 id 的 overlap class（unsharable / serial / nonextendable / shareable），包括由配置门控的 collective-serialization override。
• tracker 选择门：三个 AsyncTracker 子类（用于 TensorCore LHS 的 jellyfish TpuAsyncTracker；用于 SparseCore-offload 子 pass 的 SparseCoreAsyncTracker 和 SparseCoreResourceAwareAsyncTracker）按 pass 共存，而不是互斥选择。

Enum	xla::ResourceType（base {0..12}）由 xla::jellyfish::TpuResourceType（{13..46}）扩展
ID 数量	47（base 13 + target 34；GetNumTargetDefinedResources 返回 34）
Base 名称来源	AsyncTracker::GetResourceName @ 0x13616500（指针表 off_21920270）
Target 名称来源	TpuAsyncTracker::GetResourceName @ 0x10fff420（指针表 off_2181E148）
Op→id（base）	AsyncTracker::GetResourceTypeForOp @ 0x13612240（opcode switch）
Op→id（target）	TpuAsyncTracker::GetResourcesFromInstructionImpl @ 0x11001040 + 六个 MayAdd*
并发上限	TpuAsyncTracker::GetNumAvailableResources @ 0x10fff600
Overlap class	TpuAsyncTracker::GetResourceHazardType @ 0x110015e0（表 dword_AC0B2C0）
Tracker 安装	GetTpuAsyncTracker @ 0x10975520（TC LHS）；SC trackers 通过 SparseCoreCompiler::RunHloScheduler @ 0x1306f820
源文件	platforms/xla/service/jellyfish/latency_scheduler_cost_models_tpu.cc
置信度	CONFIRMED（字节锚定），除非行内另有说明

---

两个 Resource Enum：不要合并

	Cost-model Resource（Resource Enum）	Scheduler ResourceType（本页）
Enum	ResourceVector::Resource	ResourceType / TpuResourceType
取值	23（R[0..22]）	47（base {0..12} + target {13..46}）
度量	cycle weight（单元有多忙）	concurrency cap（多少个可共同发射）
名称来源	ResourceVectorToString @ 0x1c89bde0	GetResourceName @ 0x13616500（base）/ 0x10fff420（TPU）
按 op 映射	CycleTable::GetResource（LLO opcode → slot）	GetResourceTypeForOp / MayAdd*（HLO opcode → id）
消费者	MaxResourceCycles → bundle issue cost	GetNumAvailableResources → co-issue throttle

两个 enum 共同命名的唯一物理概念是六个 ICI ring link。cost model 把 cycle weight 放入 R[13..18]（Ici{Y,X,Z}{Plus,Minus}）；scheduler 通过 ResourceType ids 14..19（同样是 kIci{Y,X,Z}{Plus,Minus}）限制同一 link 的并发发射。这种联系是直接而非巧合的：MayAddIciLinks 读取 cost model 的六个 ICI ResourceVector 槽位，并对 cost model 存入非零 cycle count 的每个槽位发出 scheduler resource id slot+1（见下方 资源分类）。

陷阱：同名，不同数字。 若重新实现者从 cost-model 页面硬编码“ICI resources 是 13..18”，在 scheduler 中会偏一位。cost-model 槽位索引是 0xd..0x12（13..18）；scheduler ResourceType id 是 slot+1 = 0xe..0x13（14..19）。这个 off-by-one 是 MayAddIciLinks 中有意的 id = slot + 1 映射，因为 scheduler id 13 已被 kDCNbw 占用。

---

47-ID ResourceType Enum

该 enum 从两个 GetResourceName 查找函数恢复，每个函数都是由 resource id 索引的一对 {name-ptr table, length-table}。

Base XLA Resources {0..12}：AsyncTracker::GetResourceName @ 0x13616500

// AsyncTracker::GetResourceName(resource)  @ 0x13616500
const char *GetResourceName(unsigned long r) {
    if (r > 0xc) return "Not a valid default resource";   // out-of-range sentinel string
    return off_21920270[r];                               // .data.rel.ro ptr table, [rax + r*8]
}
```text

这是标准 XLA `AsyncTracker` collective 分类体系。id 是规范 XLA `ResourceType` enum 值。

| id | 名称 |
|---|---|
| 0 | `kNoResource` |
| 1 | `kAllToAll` |
| 2 | `kAllGather` |
| 3 | `kAllReduce` |
| 4 | `kCollectivePermute` |
| 5 | `kCopy` |
| 6 | `kReduceScatter` |
| 7 | `kSendRecv` |
| 8 | `kSendHost` |
| 9 | `kRecvHost` |
| 10 | `kCollectiveBroadcast` |
| 11 |（`"Not a valid default resource"` sentinel；无 enumerator） |
| 12 | `kRaggedAllToAll` |

> **注意：id 11 是一个空洞。** name-ptr table 在索引 11 没有真正的 enumerator；越界分支（`r > 0xc`）和零散索引 11 都解析到 `"Not a valid default resource"` 字符串。Ids `{2, 3, 6}`（`kAllGather`、`kAllReduce`、`kReduceScatter`）是 `SetConcurrentResourceLimits` 块从 `SchedulerConfig` 读取的三个（`xla_max_concurrent_async_all_gathers` / `all_reduces` / `reduce_scatters`）；base resources 在 target-defined availability loop 中没有条目（`AsyncTracker::GetNumAvailableResources` 返回 0）。

### Jellyfish Target Resources `{13..46}`：`TpuAsyncTracker::GetResourceName` @ `0x10fff420`

```c
// TpuAsyncTracker::GetResourceName(resource)  @ 0x10fff420  (src line 1152)
char *GetResourceName(long r) {
    CHECK(r <= 46);                                          // "resource_type < ...kTpuResourceTypeEnd"
    if (r < 13)  return AsyncTracker::GetResourceName(r);    // 0..12 → base table above
    if ((0x17FFF >> (r - 13)) & ((r - 13) < 0x11))           // 13..29, with the gap at idx 28
        return off_2181E148[r - 13];                         // TPU name table, [rax + r*8 - 0x68]
    if ((r - 30) < 0x10) return "kCustomCollective";         // 30..45 share one string
    return &nptr;                                            // 46 → empty (final catch-all)
}

GetNumTargetDefinedResources @ 0x10fff5e0 返回 34，将 target 范围固定为 [13, 13+34) = [13, 46]（含端点）。

id	名称	功能资源
13	kDCNbw	DCN（跨 slice）网络带宽
14	kIciYPlus	ICI ring link +Y
15	kIciYMinus	ICI ring link -Y
16	kIciXPlus	ICI ring link +X
17	kIciXMinus	ICI ring link -X
18	kIciZPlus	ICI ring link +Z
19	kIciZMinus	ICI ring link -Z
20	kHostToDevice	host→device DMA tap
21	kDeviceToHost	device→host DMA tap
22	kSparseCore	通用 SparseCore 引擎（按 core）
23	kSparseCoreGather	SC gather op 类
24	kSparseCoreScatter	SC scatter op 类
25	kSparseCoreDataFormatting	SC data-formatting op 类
26	kSparseCoreKernel	SC kernel op 类
27	kSparseCoreSort	SC sort op 类
28	（未命名；无 reloc，len 0：SC catch-all）	SC 通用/catch-all
29	kVmem	VMEM-resident op
30..45	kCustomCollective	16 个用户 custom-collective lane（一个共享字符串）
46	（未命名；len 0：最终 catch-all）	尾部 sentinel

注意：ids 28 和 46 匿名但有效。 off_2181E148 处的 name-ptr table 对索引 28 没有 relocation（slot 0x2181e1c0 处有空隙），id 46 落入 &nptr（空字符串）。两者都是 GetNumAvailableResources 和 GetResourceHazardType 接受的真实 resource id；只是打印路径让它们留空。Id 28 是 SparseCore “catch-all” 类别（enum1 arm），id 46 是尾部 catch-all。重新实现者必须把 resource table 大小设为 47，而不是 45。

陷阱：kIci 名称顺序不是 cost-model 槽位顺序。 GetResourceName 按 Y+, Y-, X+, X-, Z+, Z- 排列方向（ids 14..19）。cost-model IciResource 槽位标签按 X, X, Y, Y, Z, Z 排列。这两种顺序是相互独立的命名约定；运行时关系纯粹是 id = slot + 1，因此 resource id 表示的物理 link取决于上游 per-collective cost 被放入哪个 cost-model 槽位，而不是名称字符串暗示什么。不要假设 id 14 == X 的 cost slot。

---

资源分类：GetResourceTypeForOp + 六个 MayAdd* 生产者

分类分为两层。A 层（base XLA）通过单个 switch 将原始 HLO opcode 映射到 base id {0..12}。B 层（jellyfish）调用 base，然后通过六个生产者添加 target ids {13..45}，每个生产者都有自己的选择规则。

A 层：AsyncTracker::GetResourceTypeForOp @ 0x13612240

// AsyncTracker::GetResourceTypeForOp(HloOpcode op)  @ 0x13612240
long GetResourceTypeForOp(int op) {
    switch (op) {
        case 6:   return 2;    // all-gather           → kAllGather
        case 9:   return 3;    // all-reduce           → kAllReduce
        case 12:  return 1;    // all-to-all   (0xc)   → kAllToAll
        case 33:  return 10;   // collective-broadcast (0x21) → kCollectiveBroadcast
        case 34:  return 4;    // collective-permute   (0x22) → kCollectivePermute
        case 44:  return 5;    // copy                 (0x2c) → kCopy
        case 86:  return 12;   // ragged-all-to-all    (0x56) → kRaggedAllToAll
        case 93:  return 6;    // reduce-scatter       (0x5d) → kReduceScatter
        default:  return 0;    //                      → kNoResource
    }
}
```text

opcode 整数是 XLA `HloOpcode` enum 值。注意 ids 86（`ragged-all-to-all`）和 93（`reduce-scatter`）由显式 `default`-block compare 处理，而不是由覆盖 `op - 6` 的稠密 jump table 处理。

> **注意：`0x56` 是 ragged-all-to-all（id 12），`0x5d` 是 reduce-scatter（id 6）。** 很容易误读该 switch，把 key 6 绑定到 ragged-a2a；字节精确映射是 `0x56 → 12`（`kRaggedAllToAll`）和 `0x5d → 6`（`kReduceScatter`）。采用这些名称后，`SetConcurrentResourceLimits` knob→key 绑定是自洽的：key 2 ← `kAllGather`，key 3 ← `kAllReduce`，key 6 ← `kReduceScatter`。

### B 层：`TpuAsyncTracker::GetResourcesFromInstructionImpl` @ `0x11001040`

```c
// TpuAsyncTracker::GetResourcesFromInstructionImpl(hlo)  @ 0x11001040  (VLOG src 1780)
void GetResourcesFromInstructionImpl(const HloInstruction &hlo, vector<pair<id,usage>> *out) {
    AsyncTracker::GetResourcesFromInstructionImpl(hlo, out);   // (0) base ids {0..12}
    // (*) async-start/done over an all-reduce-scatter fusion → kReduceScatter (id 6)
    if ((hlo.opcode & 0xfe) == 0x10 &&                         //   async-start (0x10) | async-done (0x11)
        IsAllReduceScatterFusion(hlo.async_wrapped_instruction()))
        out->push_back({6, (this->byte208 ^ (opcode == 0x11)) + 1});
    MayAddDcnBw(hlo, out);            // (1) → id 13
    MayAddIciLinks(hlo, out);         // (2) → ids 14..19
    MayAddHostTransfers(hlo, out);    // (3) → ids 20, 21
    MayAddSparseCoreResource(hlo, out);// (4) → ids 22..27
    MayAddVmem(hlo, out);             // (5) → id 29
    MayAddCustomCollective(hlo, out); // (6) → ids 30..45
}

生产者按这个固定顺序运行。每个生产者发出 pair{resource_id, ResourceUsageType}，其中 usage 在 async-start 上是 kResourceOccupy（2 - byte208），在 async-done 上是 kResourceRelease（byte208 + 1）；byte208（this+0x208）是 start/done canonical-swap bit。

生产者	地址	发出 ids	选择规则
MayAddDcnBw	0x10fff6e0	13	跨 slice collective（opcode 位于 mask {89..111} ∩ 0x600003）；查找 CrossSliceCollectiveInfoTracker；发出 id 13
MayAddIciLinks	0x10fffb20	14..19	构建 CostModel，运行 GetCycles 写入 23 槽 ResourceVector，扫描六个 ICI 槽 {0xd..0x12}（heap table {13,14,15,16,17,18}）；对每个 cost ≠ 0 的槽发出 id slot+1
MayAddHostTransfers	0x11000280	20, 21	host send/recv → 20（H2D）/ 21（D2H）
MayAddSparseCoreResource	0x11000480	22..27	thread-name "sparsecore" gate，然后 GetSparseCoreConfig op-type enum {2..7} → ids {23,24,25,26,27}；另在 this+0x13b == 1 门控下，每个 SC core（GetNumSparseCoresUsed）发出一次 id 22
MayAddVmem	0x11000c00	29	VMEM-resident op（opcode 10/11/16/17，带 all-reduce-scatter-fusion gate）→ id 29
MayAddCustomCollective	0x11000d20	30..45	IsCustomCallAsync{Start,Done} gate；CustomCallConfig.collective_id（field 3，cfg+0x78，hasbit cfg+0x10 & 0x40）；发出 id 0x1e + collective_id，范围 [0,15]

特性：collective 按 ICI 方向重叠，而不是按一个“collective”计数器重叠。 MayAddIciLinks 不读取 opcode 来选择方向；它检查 cost model 把 cycle 放入了哪些 ICI ResourceVector 槽，并为每个槽发出匹配的 scheduler resource。两个走不同 ICI 轴的 collective（例如 +X 上的 all-reduce 和 +Y 上的 all-gather）消耗不同 resource id，可以自由重叠；同一轴上的两个会串行化。把它重新实现为单个“collective overlap”计数器会错误地串行化它们。opcode pre-filter 跳过 {6, 9, 0x22}（all-gather / all-reduce / collective-permute，它们有自己的 ring path）和 0x5d（reduce-scatter）。

特性：kCustomCollective 由数字 id 键控，而不是由 target string 键控。 MayAddCustomCollective 读取 CustomCallConfig.collective_id（int64，field 3）并计算 resource_id = 0x1e + collective_id，范围限制为 [0, 15]（越界时对 kCustomCollectiveEnd 的 CHECK fatal，消息为 "Use lower numbers of collective ids"）。因此最多 16 个不同用户 custom-collective 通过 backend config 中的显式数字 id 调度到独立资源上，custom-call target name 与资源分配无关。

---

按资源并发上限：GetNumAvailableResources @ 0x10fff600

GetNumAvailableResources(id) 返回该资源可同时在途的异步 op 数量。Base ids {0..12} 从这个循环返回 0（它们受固定的 SetConcurrentResourceLimits key block 约束，而不是 target loop）。Target ids {13..46} 读取由 TpuAsyncTracker ctor（GetTpuAsyncTracker @ 0x10975520）连线的预计算 tracker 字段。

// TpuAsyncTracker::GetNumAvailableResources(id)  @ 0x10fff600  (src line 1243)
long GetNumAvailableResources(long id) {
    CHECK(id <= 46);                                  // "...kTpuResourceTypeEnd"
    if (id < 13) return AsyncTracker::GetNumAvailableResources(id);  // base → 0
    switch (id) {
        case 13:           return this->[+0x128];     // kDCNbw
        case 20: case 21:  return this->[+0x130];     // host transfer
        case 22:           return this->[+0x140];     // kSparseCore
        case 23:           return this->[+0x148];
        case 24:           return this->[+0x150];
        case 25:           return this->[+0x158];
        case 26:           return this->[+0x160];
        case 27:           return this->[+0x168];
        case 29:           return 1;                  // kVmem — hardcoded 1
        default:                                      // 14..19, 28, 46  → [+0x170]
            if ((unsigned)(id - 30) > 0xf) return this->[+0x170];   // ici_overlap_limit
            else                          return this->[+0x178];    // 30..45 kCustomCollective
    }
}
```text

| id(s) | 名称 | tracker 字段 | available-count 来源 |
|---|---|---|---|
| 13 | `kDCNbw` | `+0x128` | `xla_tpu_dcn_overlap_limit`（`int64`，TCE `+0x11d8`） |
| 14..19 | `kIci{Y,X,Z}{±}` | `+0x170` | field 1130 `xla_tpu_sparse_core_ici_overlap_limit` |
| 20, 21 | `kHostToDevice`/`kDeviceToHost` | `+0x130` | field 803 `xla_tpu_host_transfer_overlap_limit` |
| 22 | `kSparseCore` | `+0x140` | `CoresPerChip(SC) / LogicalDevicesPerChip(SC)`（TpuTopology，按 generation） |
| 23 | `kSparseCoreGather` | `+0x148` | field 1088 `..._gather_overlap_limit` |
| 24 | `kSparseCoreScatter` | `+0x150` | field 1089 `..._scatter_overlap_limit` |
| 25 | `kSparseCoreDataFormatting` | `+0x158` | field 1090 `..._data_formatting_overlap_limit` |
| 26 | `kSparseCoreKernel` | `+0x160` | field 1091 `..._kernel_overlap_limit` |
| 27 | `kSparseCoreSort` | `+0x168` | field 1092 `..._sort_overlap_limit` |
| 28 |（SC catch-all） | `+0x170` | field 1130（与 ICI 共享） |
| 29 | `kVmem` | const | 硬编码 `1` |
| 30..45 | `kCustomCollective` | `+0x178` | 常量 `1`（ctor `push 1`） |
| 46 |（catch-all） | `+0x170` | field 1130（与 ICI 共享） |

### Field 1130：一个 knob 同时限制 ICI links *和* SC catch-all

`+0x170` 字段从 compilation-environment field **1130 = `xla_tpu_sparse_core_ici_overlap_limit`** 连线，这是一个 `AutoProto`（`AutoOr<long>`）包装器。field number 是字节精确的（`_InternalSerialize` 对 TCE `+0xa88` 的值写入 `edi = 0x46a`），切出的 `FieldDescriptorProto` 给出名称（`0x25` 字节字符串）、number（`0xea08` = 1130）、type（`TYPE_MESSAGE` → `.xla.jellyfish.AutoProto`）。在 LHS 路径中 oneof 未设置，因此 `AutoOr<long>::FromProtoOrDie` 返回 `INT64_MAX`（无上限）。

| 属性 | 值 |
|---|---|
| field number | 1130（`0x46a`） |
| field name | `xla_tpu_sparse_core_ici_overlap_limit` |
| proto type | `TYPE_MESSAGE`（`.xla.jellyfish.AutoProto`）→ `AutoOr<long>` |
| TCE `_impl_` offset | `0xa88` → tracker `+0x170`（Create arg17） |
| AUTO fallback | `INT64_MAX`（LHS 路径；无上限） |
| 被限制资源 | ids 14..19（全部六个 ICI 方向）+ 28 + 46 |

> **注意：knob 名称中的 “ici” 是字面含义，但其作用范围更宽。** `xla_tpu_sparse_core_ici_overlap_limit` 限制全部六个物理 ICI ring resource（3 个 torus 维度 × 2 个方向）*以及* SparseCore catch-all（id 28）*以及*尾部 catch-all（id 46）：这是 ICI-link-bearing 和 SC-catch-all 异步 op 的单个共享并发预算。DCN-bandwidth resource（id 13）有自己的 `int64` 上限；五个命名 SC 子类别（23..27）各有自己的 AutoProto knob。

### Id 22（`kSparseCore`）是唯一硬件派生上限

```c
// GetTpuAsyncTracker @ 0x10975520 — the id-22 (kSparseCore) count, per-gen
arg11 = 1;                                                       // default fallthrough is 1, not 0
if (EnableSparseCoreOffloadQueuingInLhs())                       // @0x1d6b81e0
    arg11 = SparseCoreOffloadQueuingOverlapLimit();              // @0x1d6b8320 — a TCE knob
else if (ShouldEnableConcurrentSparseCoreOffloading()) {        // @0x1d6b6f80
    long ldpc = Target::LogicalDevicesPerChip(kSparseCore);      // @0x1d615b00
    if (ldpc <= 0) arg11 = 0;                                    // guard against div-by-zero
    else arg11 = Target::CoresPerChip(kSparseCore) / ldpc;       // @0x1d615b40  (idiv)
}
// else: arg11 stays 1

注意：当 SparseCore offload 被禁用时，id 22 的上限是 1，不是 0。 字节精确的 GetTpuAsyncTracker @ 0x10975520 在 concurrent-offload 分支之前把默认值设为 1（v15 = 1），并且只在该分支中当 LogicalDevicesPerChip(SC) <= 0 时写入 0。因此 neither-queuing-nor-concurrent 分支保留上限 1。

Target::CoresPerChip(kSparseCore) 在 topo + coreType*0xc + 0x7c（coreType 2 = SparseCore → offset 0x94 = 148）处读取 Target[+0x3b8]（tpu::TpuTopology*，偏移 952），这是 topology 结构体中按 core-type 的 int32（TPU Topology Struct）。Target::LogicalDevicesPerChip(kSparseCore) 调用 TpuTopology::LogicalDevicesPerChip → TpuChipParts::CoreCount + TpuChipConfig::Megacore，所以除数是 megacore 折叠（在 megacore 部件上 ldpc(SC) == 2）。结果，即每 chip 物理 SC cores 除以每 chip logical devices，是唯一一个上限来自硬件数量而非配置 knob 的 target resource。

特性：id 22 的上限在 offload-queuing 下改变含义。 当设置 EnableSparseCoreOffloadQueuingInLhs（常见 embedding 生产配置）时，id 22 的上限变为 TCE knob（SparseCoreOffloadQueuingOverlapLimit），而不是 topology core-count。三路选择按字节存在；offload-queuing knob 的 field number 未解码（PARTIAL）。

---

Overlap Class：GetResourceHazardType @ 0x110015e0

两个竞争同一 resource 的 op 是否可以重叠，取决于该 resource 的 hazard class。GetResourceHazardType(id) 返回一个小整数代码。

// TpuAsyncTracker::GetResourceHazardType(id)  @ 0x110015e0  (src line 1848)
long GetResourceHazardType(long id) {
    CHECK(id <= 46);
    if (id >= 13) {
        if ((0x109FF >> (id - 13)) & ((id - 13) < 0x11))
            return dword_AC0B2C0[id - 13];                    // table for 13..29
        return 3 * (unsigned)((id - 30) >= 0x10) + 1;         // 30..45 kCustomCollective → 1; 46 → 4
    }
    // base ids 0..12:
    if (this->byte202 /*track_sync_op_resource*/ != 1)
        return AsyncTracker::GetResourceHazardType(id);       // default base: 4*(id != 5)
    // TPU collective-serialization override:
    if (id == 3 /*kAllReduce*/ || id == 6 /*kReduceScatter*/ ||
        (id == 2 /*kAllGather*/ && this->byte314))
        return 3;                                             // kSerial
    return AsyncTracker::GetResourceHazardType(id);
}
```text

hazard 代码：

| code | 含义 | overlap 行为 |
|---|---|---|
| 0 | unsharable | 单发射；该资源的任意两个 op 不重叠 |
| 1 | serial | 一个在途（FIFO 顺序） |
| 2 | nonextendable | 不能延后超过其窗口（`kVmem`） |
| 3 | serial（TPU collective override） | collective engine 单占用 |
| 4 | shareable | 可重叠到按类型的 `GetNumAvailableResources` 上限 |

对于 target ids 13..29，表为 `dword_AC0B2C0 = [0,1,1,1,1,1,1,0,0,0,0,2,0,0,0,0,2]`：

| id | 名称 | hazard | 解读 |
|---|---|---|---|
| 13 | `kDCNbw` | 0 | unsharable |
| 14..19 | `kIci{Y,X,Z}{±}` | 1 | 每方向 serial |
| 20, 21 | host transfers | 0 | unsharable |
| 22 | `kSparseCore` | 2 | nonextendable |
| 23..27 | SC sub-categories | 0 | unsharable |
| 28 |（SC catch-all） | 0 | unsharable |
| 29 | `kVmem` | 2 | nonextendable |
| 30..45 | `kCustomCollective` | 1 | serial |
| 46 |（catch-all） | 4 | shareable |

> **特性：除 `kCopy` 外，每个 base resource 都是 shareable。** base `AsyncTracker::GetResourceHazardType` 返回 `4 * (id != 5)`，也就是每个 base collective class 都是 shareable（hazard 4），*除了* `kCopy`（id 5），它是 hazard 0（unsharable）：异步 copy 在 copy engine 上串行化。TPU override 只把 `kAllReduce`/`kReduceScatter`/`kAllGather` 从 shareable 翻转为 serial（3），且仅当 `track_sync_op_resource` byte（`this+0x202`）被设置；对 all-gather 还有一个额外 byte（`this+0x314`）门控。

---

## Tracker 选择：三个 Trackers，三个子 Pass

上面的 `ResourceType` 模型是 jellyfish `TpuAsyncTracker`。它不是二进制中唯一的 `AsyncTracker`：还存在两个 SparseCore 变体，而且三者在一次编译中共存，归属于*不同*调度子 pass，而不是由单个 flag 选择。

| Tracker | Installer / call site | Resource space |
|---|---|---|
| jellyfish `TpuAsyncTracker` | `GetTpuAsyncTracker` @ `0x10975520`，来自 jellyfish `RunHloScheduler`（第 1 个 pass + field-1202 rerun） | base `{0..12}` + target `{13..46}`（本页） |
| `SparseCoreAsyncTracker` | `RunSparseCoreLatencyHidingScheduler` @ `0x1306e020` | base `AsyncTracker` + SC overrides（此处未解码） |
| `SparseCoreResourceAwareAsyncTracker` | `RunSparseCoreCostModelLatencyHidingScheduler` @ `0x1306f040`（make_shared @ `0x1306f1bb`） | 独立 `{13..17}` space，硬编码上限 |

TensorCore LHS 始终使用 jellyfish `TpuAsyncTracker`。两个 SparseCore-offload scheduler 只在 SparseCore gate 成立时运行（`SparseCoreCompiler::RunHloScheduler` @ `0x1306f820`）：

```c
// SparseCoreCompiler::RunHloScheduler gate  @ 0x1306f820
runSC =  TpuChipConfig::Megachip(Target[+0x3b8][+0x18])
      && Target::CoresPerChip(kSparseCore) > 0                  // topo[+0x94] > 0
      && (Target[+0x628] & 4  ||  Target[+0x540] != 0)          // SC-offload-enable bits
      && offloader_util::ModuleContainsLEMSparseCoreInstruction(M)
      && FLAGS_xla_sc_enable_latency_hiding_scheduler;

当 runSC 为真时，先产生 SparseCore-offload schedule：运行 RunSparseCoreLatencyHidingScheduler（plain，SparseCoreAsyncTracker），成功后 RunSparseCoreCostModelLatencyHidingScheduler（SparseCoreResourceAwareAsyncTracker + EmbeddingBackwardPassLatencyEstimator）进一步细化。然后在同一个 RunHloScheduler 中为 TensorCore LHS 安装 jellyfish TpuAsyncTracker。当 gate 失败时，SparseCore pass 回退到通用 DFSMemoryScheduler（没有 SC tracker），只有 TensorCore LHS 运行。

SparseCore-resource-aware tracker 是独立的 5-resource space

SparseCoreResourceAwareAsyncTracker 不共享 jellyfish {13..46} enum。它的 GetNumTargetDefinedResources @ 0x134a7420 返回 5，其 resources {13..17} 有自己的名称，并从 .rodata 表取得硬编码（非配置）限制：

// SparseCoreResourceAwareAsyncTracker::GetNumAvailableResources(id)  @ 0x134a7b20  (src line 261)
long GetNumAvailableResources(long id) {
    CHECK(id <= 17);                                  // "...kSparseCoreResourceTypeEnd"
    if (id < 13) return AsyncTracker::GetNumAvailableResources(id);
    return qword_AE344F8[id - 13];                    // {1, 20, 5, 1, 1}
}
```text

| id | 名称（`GetResourceName` @ `0x134a7440`） | 硬编码上限 |
|---|---|---|
| 13 | `SCS` | 1 |
| 14 | `SCT` | 20（`0x14`） |
| 15 | `ICI` | 5 |
| 16 | `LocalReduction` | 1 |
| 17 | `2DAllToAll` | 1 |

> **注意：这些是活跃的，但仅用于 SC-offload cost-model 子 pass。** `{1, 20, 5, 1, 1}` 上限可达（megachip 部件 + SC cores > 0 + offload bits + 一个 LEM instruction + flag），但它们管辖 SparseCore-offload latency-hiding pass，而*不是*主 TensorCore LHS，后者是 jellyfish `TpuAsyncTracker`。重新实现者必须让两个 resource space（`{13..46}` jellyfish 与 `{13..17}` SC-resource-aware）完全分离；它们复用相同的低整数 id 来表示完全不同的资源。

---

## 示例：两个 Collectives，一个 Schedule Step

megacore 部件上有两个独立异步 collective 的片段：

```text
%ag   = all-gather-start(%x)        ; ICI ride deposits cost into one ICI slot
%ag.d = all-gather-done(%ag)
%ar   = all-reduce-start(%y)        ; ICI ride deposits cost into a different ICI slot
%ar.d = all-reduce-done(%ar)

沿 classifier 和 cap model 走一遍：

• GetResourceTypeForOp 将 all-gather（opcode 6）→ base id 2，将 all-reduce（opcode 9）→ base id 3。
• 对 %ag，MayAddIciLinks 构建 CostModel，运行 GetCycles，并对 {0xd..0x12} 中每个 cost ≠ 0 的 ICI 槽 s 发出 resource id s + 1。假设 cost model 存入 slot 0xe → resource id 0xf = 15。对 %ar，假设它存入 slot 0xf → resource id 0x10 = 16。（某个 collective 使用哪个槽由上游 cost model 决定；resource name 顺序 Y,X,Z 独立于 slot 顺序，见上方陷阱。）
• 两个 collective 消耗不同 resource id（15 对 16）。GetResourceHazardType(15) = 1 且 GetResourceHazardType(16) = 1（每方向 serial），GetNumAvailableResources(15/16) = INT64_MAX（field 1130 AUTO），因此 scheduler 允许二者并发飞行：它们走不同 ICI ring。
• 如果二者存入同一个槽（都 → resource 15），即使上限是 INT64_MAX，serial hazard 也会强制它们按顺序执行，因为 serial resource 的两个 op 不能同时在途。
• 如果 xla_tpu_sparse_core_ici_overlap_limit（field 1130）被设置为有限 N，则最多只能有 N 个 ICI-link-bearing 异步 op（跨全部六个方向和 SC catch-all）同时 outstanding，这是共享预算。

这正是 resource model 按 ICI 方向而不是按单个“collective”类别键控 collective 的原因：物理瓶颈是按方向的 ICI link，而 cost model 已经知道每个 collective 使用哪条 link。

---

置信度摘要

Claim	Evidence
Enum 是 47 个 ID：base {0..12} + target {13..46}	GetNumTargetDefinedResources @ 0x10fff5e0 = 34；GetResourceName r <= 46 CHECK
Base 名称 {0..12}（kNoResource..kRaggedAllToAll，id 11 sentinel）	AsyncTracker::GetResourceName @ 0x13616500，表 off_21920270
Target 名称 {13..46}（kDCNbw，6× kIci，host，6× SC，kVmem，16× kCustomCollective，2 个 catch-all）	TpuAsyncTracker::GetResourceName @ 0x10fff420，表 off_2181E148
Ids 28 和 46 有效但未命名	name-ptr table 在 slot 28 有空隙；id 46 → &nptr
Op→id switch（op6→2，9→3，12→1，33→10，34→4，44→5，86→12，93→6，否则 0）	GetResourceTypeForOp @ 0x13612240
六个 MayAdd* 生产者顺序固定；usage = 通过 byte208 得到 occupy/release	GetResourcesFromInstructionImpl @ 0x11001040
MayAddIciLinks 从非零 ICI ResourceVector 槽发出 id slot+1	MayAddIciLinks @ 0x10fffb20，slot table {13..18}
MayAddCustomCollective id = 0x1e + collective_id，范围 [0,15]	MayAddCustomCollective @ 0x11000d20，CHECK kCustomCollectiveEnd
GetNumAvailableResources id→field map（+0x128..+0x178，id 29 const 1）	0x10fff600，switch byte-decoded
Field 1130 = xla_tpu_sparse_core_ici_overlap_limit 限制 ids 14..19、28、46	_InternalSerialize edi=0x46a；FieldDescriptorProto carve
Id 22（kSparseCore）上限 = CoresPerChip(SC)/LDPC(SC)（TpuTopology，按 generation）	GetTpuAsyncTracker @ 0x10975520 idiv branch
Hazard table [0,1,1,1,1,1,1,0,0,0,0,2,0,0,0,0,2]；base 4*(id!=5)；override→3	GetResourceHazardType @ 0x110015e0，dword_AC0B2C0
三个 trackers 按子 pass 共存；SC gate predicate	SparseCoreCompiler::RunHloScheduler @ 0x1306f820
SCRAAT 独立 {13..17} = SCS/SCT/ICI/LocalReduction/2DAllToAll，上限 {1,20,5,1,1}	0x134a7b20（表 qword_AE344F8）/ 0x134a7440
TCE field numbers 803/1088..1092 对应 ids 20/21/23..27	descriptor names（此处未单独字节锚定）
Id 13 DCN cap field#（507 vs 508）位于 TCE +0x11d8	int64 type + name confirmed；slot pairing not isolated
Id 22 的 offload-queuing branch field#	three-way select byte-present；knob field# not decoded

---

交叉引用

• LatencyHidingScheduler Core — 消费此 enum 的 list scheduler 和 TpuAsyncTracker dispatch；comparator 的 resource-conflict keys。
• Scheduler 概览 — LHS 在 TPU scheduling pipeline 中的位置。
• LHS ILP variant — 在 comparator 之前替换 async classifier；ResourceType model 不变。
• Resource Enum — cost-model ResourceVector::Resource（23 个槽位，cycle weight）：本页不可与之混淆的兄弟 enum。
• Bundle-Aware Cost — MaxResourceCycles bundle cost 和 MayAddIciLinks 背后的 cost model 产生的 LatencyBetween latency。
• GetHloResources Routing — cost-side resource routing，它把 ICI cycles 放入 MayAddIciLinks 读取的槽位。
• TPU Topology Struct — 其 CoresPerChip(SC)/LDPC(SC) 设置 id-22（kSparseCore）availability count 的 TpuTopology 结构体。
• Binary: extracted/libtpu-0.0.40-cp314-cp314-manylinux_2_31_x86_64/libtpu/libtpu.so（build-id 89edbbe81c5b328a958fe628a9f2207d）
• Index entry: Part VIII — Instruction Scheduling & Bundle Packing — 返回索引