GetSparseCoreConfig — Offload Op-Type 枚举来源

地址适用于 libtpu-0.0.40-cp314 wheel 中的 libtpu.so（build-id 89edbbe81c5b328a958fe628a9f2207d）。其他版本会有所不同。所有 .text/.rodata 地址均为虚拟地址；对于此二进制文件，.text VMA == file offset 0xe63c000，.rodata VMA == file offset 0x84a0000，且 .data.rel.ro VMA − 0x200000 == file offset。

摘要

GetSparseCoreConfig 是每个 SparseCore-offload 分类点都会汇入的解析器：给定一条 async HLO instruction，它解析该 instruction 的 BackendConfig，并返回一个完全物化的 xla::jellyfish::SparseCoreConfig proto（是副本，从不返回借用指针）。scheduler 关心的唯一字段是 field 2 offload，一个类型为 .xla.jellyfish.Offload 的 TYPE_ENUM，存储在 C++ 对象的 +0x24，其 presence has-bit 位于 +0x10 mask 0x4。这个枚举决定 offloaded async op 是消耗 SparseCore gather / scatter / data-formatting / kernel / sort resource lane，递归进入被包装的 collective，还是落入每 core 通用 SparseCore resource；也就是延迟隐藏 scheduler 按哪个物理 SparseCore 引擎类对该 op 进行节流。

本页恢复三个字节精确的内容。第一，GetSparseCoreConfig 解析器本身（0x1c868d20）：它的 thread-name guard（kSparseCoreThread / kSparseCoreOffloadCandidateThread）、带 globals fallback 的 SparseCoreConfig 拷贝构造，以及 consumer 端在 has-bit +0x10 & 4 后读取 +0x24 处的 offload。第二，完整的 xla::jellyfish::Offload 枚举（OFFLOAD_UNSPECIFIED 0 .. OFFLOAD_COMPUTE 8）以及每个 enumerator 精确映射到哪个 scheduler resource arm（scheduler 中为 enum − 2 索引，reservation map 中为 enum − 1）。第三，SC-offload gate bits：Target[+0x628] & 4（SC-offload-capability has-bit）和 Target[+0x540]（platform_type == 2 bool），并追踪到 jellyfish::Target::Init 设置它们的位置，以及每代默认依据（TpuVersion == 5）。

重新实现契约：

GetSparseCoreConfig(async_start) → SparseCoreConfig；consumer 只有在 has-bit +0x10 & 4 被设置时才读取 offload（field 2，对象 +0x24），否则该 op 分类为无 SC resource。
offload enumerator 选择 SC resource lane：GATHER/SCATTER/DATA_FORMATTING/KERNEL/SORT → kSparseCore{Gather,Scatter,DataFormatting,Kernel,Sort}（scheduler ids 23/24/25/26/27）；COLLECTIVE → 递归进入 async-wrapped instruction；UNSPECIFIED/EMBEDDING/COMPUTE → 无 {23..27} lane（走通用 id-22 路径）。
SC-offload scheduler sub-pass 只在 gate 成立时运行：Megachip ∧ CoresPerChip(SC) > 0 ∧ (Target[+0x628] & 4 ∨ Target[+0x540]) ∧ ModuleContainsLEMSparseCoreInstruction ∧ FLAGS_xla_sc_enable_latency_hiding_scheduler。
在真实硬件上，gate bit Target[+0x628] & 4 按代际设置；SC-offload-concurrency 默认值以 TpuVersion == 5（最新代）为 key，并可由 TCE AutoOr<bool> flag 覆盖。


解析器	`xla::jellyfish::backend_config_util::GetSparseCoreConfig` @ `0x1c868d20`
返回	`xla::jellyfish::SparseCoreConfig`（副本；ctor-copy @ `0x1d6df7c0`）
分类字段	`offload`（field 2，枚举 `.xla.jellyfish.Offload`），对象 `+0x24`，has-bit `+0x10 & 4`
Offload 枚举	9 个 enumerator，`OFFLOAD_UNSPECIFIED` 0 .. `OFFLOAD_COMPUTE` 8
Scheduler consumer	`TpuAsyncTracker::MayAddSparseCoreResource` @ `0x11000480`（索引 `enum − 2`）
Reservation consumer	`(anon)::GetSparseCoreResources` @ `0x10fdc0a0`（索引 `enum − 1`）
Gate bits	`Target[+0x628] & 4`（SC-offload-capability）∨ `Target[+0x540] != 0`（`platform_type == 2`）
Gate site	`SparseCoreCompiler::RunHloScheduler` @ `0x1306f820`
Bits set in	`jellyfish::Target::Init` @ `0x1d60fc20`
Per-gen default	`TpuVersion == 5`（`ShouldEnableConcurrentSparseCoreOffloading` @ `0x1d6b6f80`）
源文件	`platforms/xla/service/jellyfish/lowering/backend_config_util.cc`
置信度	CONFIRMED（byte-anchored），除非某行另有说明

解析器 — `GetSparseCoreConfig` @ `0x1c868d20`

解析器接收 async-start HLO，并按值返回一个 SparseCoreConfig（调用方提供 sret buffer）。它不是指针访问器：它会拷贝构造一个新的 proto，因此 consumer 的 offload/has-bit 读取发生在一个私有副本上，该副本生命周期长于 instruction 的 backend-config arena。

// backend_config_util::GetSparseCoreConfig(this_sret, async_start)  @ 0x1c868d20
SparseCoreConfig *GetSparseCoreConfig(SparseCoreConfig *this, const HloInstruction *a2) {
    if (!a2) { SparseCoreConfig::SparseCoreConfig(this, /*arena=*/0); return this; }  // null → default
    CHECK(a2->opcode() == kAsyncStart);                          // [a2+0xc] == 17  (src line 472)
    // thread-name guard (src line 475):
    auto t = a2->async_execution_thread();
    CHECK(t == "sparsecore"                                      // len 10 "sparsecore"
       || t == kSparseCoreOffloadCandidateThread);               // len 28
    auto cfg = a2->backend_config<jellyfish::BackendConfig>();    // StatusOr<BackendConfig>
    if (cfg.ok()) {
        const SparseCoreConfig *src = cfg->sparse_core_config();
        if (!src) src = &SparseCoreConfig_globals_;               // default-instance fallback
        SparseCoreConfig::SparseCoreConfig(this, /*arena=*/0, src);// copy-construct
    } else {
        SparseCoreConfig::SparseCoreConfig(this, /*arena=*/0);     // parse failed → default
    }
    /* unref the StatusOr's status rep, destroy the BackendConfig temp */
    return this;
}
```text

两个 guard 包围了解析：

- **Opcode guard。** instruction 必须是 `kAsyncStart`（opcode `17` = `0x11`）。`CHECK` 消息是 `"async_start->opcode() == HloOpcode::kAsyncStart"`（src line 472）。
- **Thread-name guard。** async-execution thread 必须是 `"sparsecore"`（10 字节的 `kSparseCoreThread`，通过 immediate `0x6F63657372617073` = `"sparsec"` + `0x6572` = `"or"` 低字节比较）**或** 28 字节的 `kSparseCoreOffloadCandidateThread`。这是反汇编中的 SIMD `vptest` compare；`CHECK` 消息列出两个 thread（src line 475）。

> **注意 — 解析器返回 default，而不是 null pointer。** 对于 null instruction、非 `kAsyncStart` opcode（在 `MayAdd*` 调用方中会预先过滤，而不是由这里 `CHECK`）、解析失败、或 `BackendConfig` 不含 `sparse_core_config` 子消息，结果都是 default-constructed `SparseCoreConfig`，所有 has-bit 清零。因此 consumer 的 `has_offload()` 测试（`+0x10 & 4`）才是真正的 gate：default proto 的 bit 清零，并分类为“无 SC resource”。重新实现不能把该解析器当成调用点上的 fallible 函数；其 fallibility 已折叠到清零的 has-bit 中。

---

## `SparseCoreConfig` Proto — 字段图

`GetSparseCoreConfig` 返回完整的 `SparseCoreConfig`；scheduler/reservation 分类器只读取 `offload`（field 2），但完整字段图（来自 `SparseCoreConfig::_InternalSerialize` @ `0x1d6dfae0`）固定了对象布局，并确认 `offload` 位于对象 `+0x24`，由 has-bit `+0x10 & 4` 保护。

| field | name | object off | has-bit（`obj+0x10`） | proto type |
|---|---|---|---|---|
| 1 | `tiling` | `+0x20`（i32） | `0x2` | enum `.xla.jellyfish.Tiling` |
| 2 | **`offload`** ← SC op-type 枚举 | `+0x24`（i32） | `0x4` | enum `.xla.jellyfish.Offload` |
| 3 | `comp_env` | `+0x18`（msg） | `0x1` | msg `ScCompilationEnvironment` |
| 4 | `enable_megacore` | `+0x2c`（bool） | `0x10` | bool |
| 5 | `hbm_bandwidth_adjustment_factor` | `+0x28`（f32） | `0x8` | float（fixed32） |
| 6 | `function_mode` | `+0x2d`（byte） | `0x20` | enum/bool |
| 7 | `dedup_id` | `+0x30`（i64） | `0x80` | int64 |
| 8 | `enable_program_barrier` | `+0x2e`（bool） | `0x40` | bool |
| 9 | `load_dat` | `+0x38`（bool） | `0x100` | bool |

`offload` 的 `FieldDescriptorProto` 已按字节切出：tag `0a07`（`"offload"`，7 字节）、`1802`（number = 2）、`2001`（label = optional）、`280e`（type = `TYPE_ENUM` = 14）、`3216`（`.xla.jellyfish.Offload`，0x16 字节）。描述符类型字符串 `.xla.jellyfish.Offload`、`.xla.jellyfish.Tiling`，以及字段名 `enable_megacore` / `hbm_bandwidth_adjustment_factor` / `function_mode` / `dedup_id` / `enable_program_barrier` / `load_dat` 都存在于 `.rodata`。

> **陷阱 — `offload` 是 backend-config 枚举，不是 custom-call target 或 MLIR op kind。** 由 `GetSparseCoreConfig` 设 key 的 op-type 分类是 *instruction backend config 上的 proto 字段*。这与普通 `SparseCoreAsyncTracker` 不同，后者通过单独的枚举（`SparseCoreOperationType`）按 custom-call *target string*（`"AllToAllDynamic"`）设 key；见 [Plain Tracker 按 Target Name 设 Key，而不是 `offload`](#plain-tracker-按-target-name-设-key而不是-offload)。不要混淆两者：`offload`（本页）和 `SparseCoreOperationType` 是独立枚举，拥有独立的值空间和独立的 consumer。

---

## `xla::jellyfish::Offload` 枚举

从 `EnumDescriptorProto` 精确按字节解码（每个值都是 `10 NN` `EnumValueDescriptorProto`）。全部九个 enumerator 名称字符串（`OFFLOAD_UNSPECIFIED` .. `OFFLOAD_COMPUTE`）都存在于 `.rodata`。

| value | enumerator | 语义 |
|---|---|---|
| 0 | `OFFLOAD_UNSPECIFIED` | unset / default |
| 1 | `OFFLOAD_EMBEDDING` | embedding lookup/update offload |
| 2 | `OFFLOAD_GATHER` | gather op-class |
| 3 | `OFFLOAD_SCATTER` | scatter op-class |
| 4 | `OFFLOAD_COLLECTIVE` | collective（递归进入 wrapped op） |
| 5 | `OFFLOAD_DATA_FORMATTING` | data-formatting op-class |
| 6 | `OFFLOAD_KERNEL` | 通用 SC kernel op-class |
| 7 | `OFFLOAD_SORT` | sort op-class |
| 8 | `OFFLOAD_COMPUTE` | compute offload |

---

## Scheduler Consumer — `MayAddSparseCoreResource` @ `0x11000480`

`TpuAsyncTracker::MayAddSparseCoreResource` 是把 offloaded async op 转换为 scheduler resource id 的 producer。它先重新应用同一 thread guard（async + thread `"sparsecore"` / `kSparseCoreOffloadCandidateThread`，否则返回），调用 `GetSparseCoreConfig`，然后**仅在 has-bit 被设置时**对 `offload` 进行 switch。switch index 直接是 `offload`（`add 0xfffffffe` / jump table 覆盖值 `{2..7}`，即 table base 后的 `enum − 2`）。

```c
// MayAddSparseCoreResource(this, async_start, &out)  @ 0x11000480  (condensed)
SparseCoreConfig cfg; GetSparseCoreConfig(&cfg, async_chain_start(async_start));
int usage = /* 2 = kResourceOccupy on start, 1 = kResourceRelease on done */;
if (cfg.has_bits & 0x4) {                       // v55 & 4  → has_offload()
    switch (cfg.offload) {                      // v56 = [&cfg + 0x24]
        case 2: out.push_back({23, usage}); break;  // OFFLOAD_GATHER          → kSparseCoreGather
        case 3: out.push_back({24, usage}); break;  // OFFLOAD_SCATTER         → kSparseCoreScatter
        case 4: MayAddSparseCoreResource$_0(&out,    // OFFLOAD_COLLECTIVE      → recurse on
                  async_wrapped_instruction(async_start), usage); break;  //    async-wrapped op
        case 5: out.push_back({25, usage}); break;  // OFFLOAD_DATA_FORMATTING → kSparseCoreDataFormatting
        case 6: out.push_back({26, usage}); break;  // OFFLOAD_KERNEL          → kSparseCoreKernel
        case 7: out.push_back({27, usage}); break;  // OFFLOAD_SORT            → kSparseCoreSort
        default: break;                             // 0/1/8 → no {23..27} arm
    }
}
// independent of the offload switch — the general per-core SparseCore resource (id 22):
if (this[+0x13b] == 1) {                         // a1+315 — "per-core" gate
    for (i = 0; i < GetNumSparseCoresUsed(async_chain_start(async_start)); i++)
        out.push_back({22, usage});              // id 22 once per used SC core
} else {
    out.push_back({22, usage});                  // id 22 once
}

`offload`	enumerator	scheduler arm（`MayAddSparseCoreResource`，idx `enum − 2`）
0	`OFFLOAD_UNSPECIFIED`	无 arm — 仅 id 22 路径
1	`OFFLOAD_EMBEDDING`	无 arm — 仅 id 22 路径
2	`OFFLOAD_GATHER`	id 23 `kSparseCoreGather`
3	`OFFLOAD_SCATTER`	id 24 `kSparseCoreScatter`
4	`OFFLOAD_COLLECTIVE`	递归进入 `async_wrapped_instruction`（`$_0` @ `0x110008e0`）
5	`OFFLOAD_DATA_FORMATTING`	id 25 `kSparseCoreDataFormatting`
6	`OFFLOAD_KERNEL`	id 26 `kSparseCoreKernel`
7	`OFFLOAD_SORT`	id 27 `kSparseCoreSort`
8	`OFFLOAD_COMPUTE`	无 arm — 仅 id 22 路径

resource id {22..27} 是 47-id scheduler ResourceType 枚举中的 SparseCore 引擎类；其 concurrency cap 与 hazard class 见 ResourceType 分类。

怪异点 — OFFLOAD_COLLECTIVE 会递归，而不是发射 lane。 值 4 是唯一不 push kSparseCore* id 的 arm。它会在 async-wrapped instruction 上重新进入 MayAddSparseCoreResource（0x110008e0 处的 $_0 lambda），因此包装在 SC-offload async op 中的 collective 会由被包装 op 自己的 offload 字段分类。把值 4 映射为 resource id 的重新实现会重复计数或错误节流 wrapped collective。
注意 — id-22 通用路径独立于 offload。 per-core kSparseCore（id 22）发射发生在 offload switch 之后，并由单独的 byte（this+0x13b == 1）设门，选择“每个已用 SC core 一个 id-22”（GetNumSparseCoresUsed）还是“一个 id-22”。因此即使是没有命中 {23..27} arm 的 OFFLOAD_UNSPECIFIED/EMBEDDING/COMPUTE op，也仍会消耗通用 SparseCore resource。offload 枚举细化的是哪个子引擎；id 22 是始终存在的每 core occupancy。

Reservation Consumer — `GetSparseCoreResources` @ `0x10fdc0a0`

reservation-map 对偶读取同一个 SparseCoreConfig.offload 字段（同一个 +0x24 值，同一个 +0x10 & 4 has-bit），但用 enum − 1 索引（一次 dec 后 cmp 6），因此 live range 是 {1..7}：它额外覆盖了拥有自己 arm 的 OFFLOAD_EMBEDDING（值 1）。

// (anon)::GetSparseCoreResources(async_start)  @ 0x10fdc0a0  (condensed)
SparseCoreConfig cfg; GetSparseCoreConfig(&cfg, async_start);
if ((cfg.has_bits & 4) != 0) {                  // v29[16] & 4  → has_offload()
    switch (cfg.offload) {                      // index = enum, table base enum−1, range {1..7}
        case 1: /* OFFLOAD_EMBEDDING → embedding/general reservation arm */ ...
        case 2: /* OFFLOAD_GATHER          → kSparseCoreGather   */ ...
        case 3: /* OFFLOAD_SCATTER         → kSparseCoreScatter  */ ...
        case 4: /* OFFLOAD_COLLECTIVE      → collective arm      */ ...
        case 5: /* OFFLOAD_DATA_FORMATTING → kSparseCoreDataFmt  */ ...
        case 6: /* OFFLOAD_KERNEL          → kSparseCoreKernel   */ ...
        case 7: /* OFFLOAD_SORT            → kSparseCoreSort     */ ...
    }
}
```text

> **陷阱 — scheduler 用 `enum − 2` 索引，reservation 用 `enum − 1`。** 两个 consumer 读取同一个 proto 字段、同一个偏移，但 jump-table base 差一。scheduler（`MayAddSparseCoreResource`）的 dense table 从值 2（`OFFLOAD_GATHER`）开始，所以 `OFFLOAD_EMBEDDING`（1）没有命中 scheduler arm；reservation map（`GetSparseCoreResources`）从值 1 开始，所以会为 embedding 做 reservation。重新实现者必须保留两个 index base：不要假设两个分类器共享同一个 switch。

reservation-map 的 resource→limit 结构见 [SC Queue Assignment & Reservation](sc-queue-assignment-reservation.md)。

---

## SC-Offload Gate Bits

SparseCore-offload scheduler sub-pass 只有在 `SparseCoreCompiler::RunHloScheduler`（@ `0x1306f820`）发现 gate 谓词为真时才运行。其合取项中有两个 `Target` bitfield 读取：

```c
// SparseCoreCompiler::RunHloScheduler gate  @ 0x1306f820  (object offsets in bytes)
runSC =  TpuChipConfig::Megachip( Target[+0x3b8][+0x18] )                 // @0x1306f84c
      && *(int*)( Target[+0x3b8] + 0x94 ) > 0       // CoresPerChip(kSparseCore) > 0  @0x1306f863
      && ( (Target[+0x628] & 4) != 0  ||  Target[+0x540] != 0 )           // the two gate bits
      && offloader_util::ModuleContainsLEMSparseCoreInstruction(M)        // @0x1306fbc8
      && FLAGS_xla_sc_enable_latency_hiding_scheduler;                    // @0x1306fc04

反编译中这是 (*((_BYTE*)this + 1576) & 4) != 0 || *((_BYTE*)this + 1344)：byte 1576 = 0x628，byte 1344 = 0x540。*(int*)(... + 148) > 0 是 0x94 CoresPerChip(SC) 读取。整个谓词出现两次（eager check 与 SC-path re-test）。

bits 的设置位置 — `Target::Init` @ `0x1d60fc20`

两个字段都在 jellyfish::Target::Init 中写入。寄存器 r12/v342 是被初始化的 Target*；v98/*v98 是 TpuTopology 的第一个标量（platform-type enum）。

// jellyfish::Target::Init  @ 0x1d60fc20  (relevant writes)
Target[+0x540] = (TpuTopology[+0] == 2);     // platform_type == 2   (decompile: _R12+1344)
Target[+0x541] = (TpuTopology[+0] == 1);     // platform_type == 1   (decompile: _R12+1345)
// inside the predicate-gated config-append loop:
Target[+0x628] |= 1;                         // bit-0  (config sub-field A has-bit)  @0x1d611d52
Target[+0x628] |= 4;                         // bit-2  (SC-offload-capability has-bit) @0x1d612121
```text

- **`Target[+0x540]`** 是一个 `bool` = `(TpuTopology[+0] == 2)`。第一个 `TpuTopology` 标量是内部 platform-type enum；值 `2` 是 `iss`（simulator）路径（`platform_type == 1` 落入相邻 byte `Target[+0x541]`）。因此无论 capability bit 如何，`Target[+0x540] != 0` 都会在 simulator 上强制进入 SC 路径。
- **`Target[+0x628]`** 是 `_has_bits_` 风格的 qword（反编译为 `*((_QWORD*)_R12 + 197)`，qword index 197 = byte `0x628`）。bit-2（mask `0x4`）在一个展开的 config-append loop 内 OR 进去，而该 loop 自身由 SC-offload feature-detect 设门；它是**SC-offload-capability has-bit**，为符合条件的（最新代）部件设置。bit-0（mask `0x1`）在同一个 loop 中更早被 OR，用于相邻 config sub-field。gate 谓词 `(Target[+0x628] & 4) == 0 → read Target[+0x540]` 在 `Target::Init` 自身中也逐字重放（并与 `Megachip ∧ CoresPerChip(SC) > 0` 组合），也就是 SC-offload feature-detect。

| gate bit | object off | 含义 | set in `Target::Init` |
|---|---|---|---|
| `Target[+0x628] & 4`（bit-2） | `+0x628` qword | SC-offload-capability has-bit（按代际、由谓词设门） | OR'd `\|= 4` @ `0x1d612121` |
| `Target[+0x540]` | `+0x540` byte | `platform_type == 2`（iss/simulator） | `= (TpuTopology[+0] == 2)` @ `0x1d610b1b` |
| `Target[+0x541]` | `+0x541` byte | `platform_type == 1`（相邻字段，不在 gate 中） | `= (TpuTopology[+0] == 1)` @ `0x1d610b29` |
| `Target[+0x628] & 1`（bit-0） | `+0x628` qword | 相邻 config-append has-bit（不在 gate 中） | OR'd `\|= 1` @ `0x1d611d52` |

> **注意 — bit-2 的精确 proto sub-field 名称没有被隔离出来。** bit-set site、mask 与 value 都是字节精确的，并且该 bit 位于 predicate-gated config-append loop 中，与 bit-0 和两个 SSO string（对象 `+0x580`/`+0x5f0`）相邻。但命名此 SC-offload-capability sub-field 的单个 descriptor entry（一个从芯片 `vector_isa`/`TpuSequencerParts` 拷贝出来的嵌套 config 字段）没有钉到某个描述符上。无论如何，其*角色*——scheduler gate 读取的 SC-offload-capability flag——是字节精确的。**置信度：bit position CONFIRMED；sub-field proto name INFERRED。**
>
> **陷阱 — `platform_type` 枚举顺序来自 descriptor-string order。** `TpuTopology[+0]` 是 topology 的第一个标量 = 内部 `TpuPlatformType` 枚举（依据 `ValidateArgs(TpuPlatformType, …)` 签名以及已证明的 `TpuTopology[+0x8] = TpuChipParts*` 布局）。gate 将其与 `== 2` 比较；value→name 配对 `{0 hardware, 1 grm, 2 iss}` 取自 descriptor-string order（`TpuPlatformTypeToProto = type + 1`），不是单独解码的 `platform_type()` getter。**置信度：`== 2` 比较及其 gate 角色 CONFIRMED；enum value→name 配对 INFERRED。**

---

## 每代默认依据 — `TpuVersion == 5`

输入 scheduler 的两个 SC-offload-concurrency knob 在恰好一个芯片代际上默认*启用*。两者都会把 hardware default 计算为 `TpuChipParts[+0] == 5`（即 `tpu::TpuVersion == 5`），然后允许 `TpuCompilationEnvironment` `AutoOr<bool>` flag 通过 `0x100` “is-set” bit 覆盖。

```c
// ShouldEnableConcurrentSparseCoreOffloading(tce_view, topo, b)  @ 0x1d6b6f80
hw_default = (TpuChipParts[+0] == 5) & ~b;          // *(_DWORD*)(topo+8) == 5  → TpuVersion 5
flag = tce[+0x458] ? tce[+0x458] : &AutoProto_globals_;   // a1 + 1112 = 0x458
v = AutoOr<bool>::FromProtoOrDie(flag);
return (v & 0x100) ? /*flag set → use flag value*/ : hw_default;

// EnableSparseCoreOffloadQueuingInLhs(tce_view, topo)  @ 0x1d6b81e0
hw_default = (TpuChipParts[+0] == 5);               // *(_DWORD*)(topo+8) == 5
flag = tce[+0x730] ? tce[+0x730] : &AutoProto_globals_;   // a1 + 1840 = 0x730
... same AutoOr<bool> 0x100 override ...

TpuChipParts[+0] 是 TpuVersion（0-based 芯片代际枚举）；TpuChipParts::ToProto → TpuVersionToProto(v) = v + 1（已确认：反编译体就是 return v + 1）。来自 TpuVersionFromString（@ 0x20b3a5a0，init-list @ 0x220117b0）的 codename 表：

TpuVersion（internal）	codename	proto value（`= internal + 1`）
0	`jellyfish`	1
1	`dragonfish`	2
2	`pufferfish`	3
3	`viperfish`	4
4	`ghostlite`	5
5	`6acc60406`	6 ← SC-offload concurrency default ON

六个 codename 都逐字存在于 .rodata（init-list off_220117B0；value-5 字符串 "6acc60406" 位于 .rodata VMA 0x863f0cf，len 9，与 TpuVersion == 5 配对）。完整 generation map 见 TPU Version Codename Matrix。

注意 — override flag 的 field number 未解码。 两个 knob 都在 TCE _impl_ offset 0x458（concurrency）/ 0x730（offload-queuing-in-LHS）读取一个 AutoOr<bool>，unset 时回退到 AutoProto_globals_；offset 和 0x100-bit override 是字节精确的，但两个 proto field number（这些 offset 对应的 _InternalSerialize tag）未被隔离。置信度：offset + override mechanism CONFIRMED；field number PARTIAL。 见 TpuCompilationEnvironment 和 TCE Field Offsets & Defaults。

Plain Tracker 按 Target Name 设 Key，而不是 `offload`

并非每个 SparseCore scheduler 都以 GetSparseCoreConfig 为 key。当 gate 成立时，SC-offload schedule 会先由一个普通 SparseCoreAsyncTracker 生成，而该 tracker 按 opcode + custom-call target name 对 async-schedulable op 分类，而不是按 offload backend enum。这是看清二进制文件中 offload 只是两套独立 SC 分类机制之一的最清楚方式。

// SparseCoreAsyncTracker::IsSupportedAsyncStart(hlo)  @ 0x134964c0
bool IsSupportedAsyncStart(const HloInstruction *h) {
    int op = h->opcode();                          // [h+0xc]
    if (op == 12) return true;                     // 0xc  all-to-all
    if (op == 17) return true;                     // 0x11 async-start
    if (op == 49)                                  // 0x31 custom-call
        return SparseCoreOperationTypeFromString(h->custom_call_target()) == 8;  // "AllToAllDynamic"
    return false;
}
// IsSupportedAsyncDone @ 0x13496520 is identical except op 16 (0x10 async-done) replaces 17.
```text

`SparseCoreOperationTypeFromString`（@ `0x14b7f060`）是一个在*另一个*枚举 `SparseCoreOperationType` 上进行 chained `EqualsIgnoreCase` 映射的函数；其前八个值已按顺序确认：`1 SparseMap`、`2 CooToCsr`、`3 CooToEll`、`4 SparseMapRow`、`5 SortLexicographic`、`6 ReduceDuplicates`、`7 EllToCsr`、`8 AllToAllDynamic`（之后还有 `9 ScSendToTc`、`10 ScReceiveFromTc` 等，远超 8）。plain tracker 只对 `== 8`（`"AllToAllDynamic"`）设门；它覆盖 SC all-to-all op，并且（通过 `PostProcessScheduleGraph` → `FindNearestAllToAlls` @ `0x13496600`）把 schedule 偏向这些 op。

| classifier | keyed on | enum | values used |
|---|---|---|---|
| `MayAddSparseCoreResource`（本页） | `SparseCoreConfig.offload`（field 2） | `xla::jellyfish::Offload` | `{2..7}` → ids `{23..27}` + recurse |
| `GetSparseCoreResources`（reservation） | `SparseCoreConfig.offload`（field 2） | `xla::jellyfish::Offload` | `{1..7}` |
| `SparseCoreAsyncTracker::IsSupportedAsync{Start,Done}` | custom-call target string | `SparseCoreOperationType` | `== 8`（`AllToAllDynamic`） |

> **怪异点 — 两个枚举，一个词组 "SparseCore op type"。** `xla::jellyfish::Offload`（backend-config 枚举，9 个值，驱动 resource lane）与 `SparseCoreOperationType`（custom-call target-name 枚举，≥ 8 个值，驱动 async-schedulability）很容易混淆，因为两者都描述“这是哪类 SparseCore op”。它们完全独立：不同 namespace、不同值空间、不同读取路径、不同 consumer。`GetSparseCoreConfig` 解析前者；`custom_call_target()` + `SparseCoreOperationTypeFromString` 解析后者。

---

## 工作示例 — 最新代部件上的 `OFFLOAD_SCATTER` Async Op

一个 SparseCore-offloaded scatter，运行在 megacore 最新代（`TpuVersion == 5`）部件上：

```text
%sc = async-start(%coo), execution_thread="sparsecore",
        backend_config = { sparse_core_config { offload: OFFLOAD_SCATTER } }
%sc.d = async-done(%sc)

遍历解析器和分类器：

Gate。 Target::Init 为最新代部件设置了 Target[+0x628] |= 4，因此 (Target[+0x628] & 4) != 0 为真；在 Megachip、CoresPerChip(SC) > 0、module 中存在 LEM SparseCore instruction 且 LHS flag 开启时，SparseCoreCompiler::RunHloScheduler 进入 SC-offload sub-pass。
解析。 对 %sc，MayAddSparseCoreResource 重新检查 thread（"sparsecore"，len 10，通过），调用 GetSparseCoreConfig，后者拷贝构造 SparseCoreConfig。has_offload()（+0x10 & 4）已设置；offload（+0x24）读取为 3（OFFLOAD_SCATTER）。
分类。 对 3 的 switch 在 async-start 上发射 {24, kResourceOccupy}，也就是 scheduler resource id 24（kSparseCoreScatter），并在 %sc.d 上发射匹配的 {24, kResourceRelease}。独立地，id-22 路径发射通用 per-core kSparseCore occupancy。
节流。 scheduler 把并发 id-24 op 限制在 GetNumAvailableResources(24)（..._scatter_overlap_limit TCE knob），并把 id 24 视为不可共享（hazard 0）；见 ResourceType 分类。
Default-instance 对比。 如果 backend config 省略了 sparse_core_config，GetSparseCoreConfig 会返回 default instance，has_offload() 清零，%sc 只会消耗通用 id-22 lane，不会有 scatter-specific throttle。

这正是为什么 offload 是 SC op-type 分类器：它是把一个 generic SparseCore async op 细化成 scheduler 可建模的具体 engine-class resource 的唯一 proto 字段。

置信度摘要

论断	证据
`GetSparseCoreConfig` 返回 copy-constructed `SparseCoreConfig`（globals fallback）	`0x1c868d20`：`SparseCoreConfig::SparseCoreConfig(this, 0, src)`、`SparseCoreConfig_globals_`
Thread guard：`"sparsecore"`（len 10）∨ `kSparseCoreOffloadCandidateThread`（len 28）；opcode `kAsyncStart`	`0x1c868d20` SIMD compare + `CHECK` strings（src 472/475）
`offload` = field 2，枚举 `.xla.jellyfish.Offload`，对象 `+0x24`，has-bit `+0x10 & 4`	`_InternalSerialize` @ `0x1d6dfae0`；`FieldDescriptorProto` carve；consumer reads
Offload enum `OFFLOAD_UNSPECIFIED` 0 .. `OFFLOAD_COMPUTE` 8（9 个值）	`EnumDescriptorProto` @ `0xbfa1f9f`；全部 9 个字符串在 `.rodata`
Scheduler arm map `{2→23, 3→24, 4→recurse, 5→25, 6→26, 7→27}`，idx `enum − 2`	`MayAddSparseCoreResource` @ `0x11000480` switch（反编译）
Reservation map 读取同一字段，idx `enum − 1`，范围 `{1..7}`（覆盖 `EMBEDDING`）	`GetSparseCoreResources` @ `0x10fdc0a0`：`v29[16] & 4`，`case 1..7`
id-22（`kSparseCore`）路径独立于 `offload`，由 `this+0x13b == 1` 设门	`0x11000480`：post-switch loop on `GetNumSparseCoresUsed`
Gate predicate `(Target[+0x628] & 4) ∨ Target[+0x540]` 加 Megachip/CoresPerChip(SC)/LEM/flag	`RunHloScheduler` @ `0x1306f820`：`this+1576 & 4`、`this+1344`、`+148 > 0`
`Target[+0x540] = (TpuTopology[+0] == 2)`；`Target[+0x541] = (== 1)`	`Target::Init` @ `0x1d60fc20` lines `_R12+1344/+1345 = *v98 == 2/1`
`Target[+0x628] \|= 4`（bit-2）与 `\|= 1`（bit-0）位于 predicate-gated config loop	`Target::Init`：`v342+197 = … \| 4` / `\| 1`（qword 197 = `0x628`）
命名 bit-2 的 proto sub-field（SC-offload-capability）	bit-set site 字节精确；descriptor entry 未隔离
`platform_type` value→name 配对 `{0 hardware, 1 grm, 2 iss}`	descriptor-string order + `ToProto = type+1`；`== 2` 比较字节精确
Per-gen default `TpuVersion == 5`；TCE `+0x458`/`+0x730` 处 `AutoOr<bool>` 0x100 override	`0x1d6b6f80`（`*(topo+8)==5`、`a1+1112`）；`0x1d6b81e0`（`a1+1840`）
两个 SC-offload knob 的 TCE override field number	offset + override 字节精确；field number 未解码
TpuVersion 0..5 = jellyfish/dragonfish/pufferfish/viperfish/ghostlite/`6acc60406`；proto = +1	`TpuVersionFromString` init-list @ `0x220117b0`（v5 字符串 `"6acc60406"` @ `.rodata 0x863f0cf`）；`TpuVersionToProto` body `v+1`
Plain `SparseCoreAsyncTracker` 按 opcode `{0xc, 0x11/0x10, 0x31}` + `SparseCoreOperationType == 8` 设 key	`IsSupportedAsyncStart/Done` @ `0x134964c0`/`0x13496520`；`FromString` @ `0x14b7f060`
`SparseCoreOperationType` 值 1..8 = SparseMap..AllToAllDynamic	`0x14b7f060` chained `EqualsIgnoreCase`（反编译，按顺序）

交叉引用

SparseCore 概览 — offloaded op-class 所面向的 SCS/TAC/TEC 引擎模型。
SparseCore 架构 — kSparseCore{Gather,Scatter,DataFormatting,Kernel,Sort} resource lane 背后的硬件引擎。
SC 后端 Pipeline — SC-offload scheduling 在 SparseCore compile flow 中的位置。
Stream Gather/Scatter — OFFLOAD_GATHER / OFFLOAD_SCATTER 路由到的 gather/scatter datapath。
SC Queue Assignment & Reservation — 本分类器的 reservation-map（GetSparseCoreResources）对偶。
SparseCore vs Neuron MatMultSparse — sparse offload model 的跨架构对比。
ResourceType 分类 — 47-id scheduler 枚举；本页发射的 ids 22..27 的 cap/hazard，以及三 tracker 选择 gate。
LatencyHidingScheduler Core — 消费 resource id 的 list scheduler。
TpuCompilationEnvironment — 携带 AutoOr<bool> SC-offload override flag（+0x458 / +0x730）的 TCE。
TCE Field Offsets & Defaults — TCE _impl_ offset 表。
TPU Version Codename Matrix — TpuVersion 0..5 generation map（default-on basis = version 5）。
TPU Topology Struct — 其第一个标量（platform_type）设置 Target[+0x540] 的 TpuTopology。
Binary: extracted/libtpu-0.0.40-cp314-cp314-manylinux_2_31_x86_64/libtpu/libtpu.so（build-id 89edbbe81c5b328a958fe628a9f2207d）
索引条目: Part IX — SparseCore & BarnaCore / SparseCore back-end — 返回索引

GetSparseCoreConfig — Offload Op-Type 枚举来源 ​

摘要 ​

解析器 — GetSparseCoreConfig @ 0x1c868d20 ​

Reservation Consumer — GetSparseCoreResources @ 0x10fdc0a0 ​

bits 的设置位置 — Target::Init @ 0x1d60fc20 ​

Plain Tracker 按 Target Name 设 Key，而不是 offload ​

置信度摘要 ​

交叉引用 ​