TensorCore 屏障分配和 `InitializeOnScs`

此页上的每个地址、字段偏移、操作码值和枚举值都是从 libtpu-0.0.40-cp314 轮中的 libtpu.so 读取的(构建 ID 89edbbe81c5b328a958fe628a9f2207d；构建 libtpu_lts_20260413_b_RC00；781,691,048 B，未剥离 — 完整 C++ 符号)。 .text VMA 等于 0xe63c000 处的文件偏移量； .rodata / .lrodata 是身份映射的。所有地址都是VMA。其他车轮版本有所不同。

摘要

该页面拥有两个字节锚定部分，它们包围 TensorCore (TC) 片上屏障路径。第一个是 TC 障碍分配通道 — TensorCoreBarrierAssignment::Run @0x109c7420 及其每键类型选择器 DetermineBarrierConfigForKey @0x109c6fa0 — SparseCore 着色/分配的 TensorCore 对应项。它遍历每个 TC 集体，为每个操作构建一个 TensorCoreBarrierKey，提供两个贪婪着色通道的冲突设置，并为每个不同的键分配一个 BarrierConfig {type, id}，该 BarrierConfig {type, id} 被写回到集体的 HLO BackendConfig 中。第二个是 ExplicitUniDirRingStrategy::InitializeOnScs @0x1337aa60 及其策略+0x98 查找回调(ExplicitRingRecord @0x133a9a40 用于普通环，ExplicitAllToAllRingRecord @0x133a94a0 用于全对所有双胞胎)： SparseCore 端运行时 init 将静态显式环配置表转换为已编程的 next_chip / ordinal / reorder 三重绑定到策略对象上。

这是由芯片 SFLAG 模型连接的两个不同的子系统。 TC 分配过程决定每个键的屏障类型(全局与每个键)，而 InitializeOnScs 回调绑定SC 定序器的屏障 sync_add 目标的每个环传输几何结构 - 两者最终都对同一芯片 SFLAG 块进行编程。通用贪婪着色引擎归屏障着色所有(该页面消耗其冲突集，不会重新派生)； BarrierConfig.id → 具体的 SFLAG 数字公式归屏障 → SFLAG 绑定所有(此页面生成 {type, id}，它不会降低它)。 BarrierType 枚举、InferBarrierConfig 标准化器和保留块几何位于概述上。

重新实现，合约为：

TC 按图形颜色进行删除，SC 按静态密钥哈希进行删除。 Run 运行两个 BarrierColoring 通道 (屏障着色)，合并它们的冲突集，然后每个不同的 TensorCoreBarrierKey 调用 DetermineBarrierConfigForKey(key, config, has_conflict)。非冲突密钥可以共享 REPLICA(2) 或采用 GLOBAL(1) 屏障；冲突的密钥(在合并的冲突集中)被强制为新的 CUSTOM(3) id。
DetermineBarrierConfigForKey 仅写入 {1,2,3}。 键入消息 +0x20，id 为 +0x18，存在状态为 +0x10 |= 3。 GLOBAL(1) → id = -1； REPLICA(2) / CUSTOM(3) → id = replica_group_count (key+0x10)，当 std::map<key, BarrierConfig> 去重发现密钥已存在时选择 2。 MEGACORE(4) 从未写在这里。
IsGlobalBarrierBeneficial 是一种狭窄的启发式方法。 Global 仅对于通道 all-to-all(0xc) 是“有益的”，其测试的拓扑轴是单位大小的并且跨越单个副本组 - 自定义每键屏障将同步一个平凡组的退化情况。
InitializeOnScs 是静态配置 → 硬件最终跳跃。 它通过 CoresPerChip / LogicalDevicesPerChip 折叠运行时核心索引，然后使用捕获的表描述符 (strategy+0x88) 和运行时 global_core_id / ordinal 调用 [strategy+0x98]；回调索引 const-literals 表并返回 (next_chip, ordinal, reorder)，写入 strategy+0x58 (ordinal_) / +0x60 (next_chip_) / +0x78 (reordering_map_)。 next_chip_ 是每环 SyncAdd 的目标。


TC分配通行证	`xla::jellyfish::TensorCoreBarrierAssignment::Run` @`0x109c7420`
每键类型选择器	`DetermineBarrierConfigForKey` @`0x109c6fa0`(写入 `{1,2,3}`；从不写入 `4`)
全局启发式	`IsGlobalBarrierBeneficial` @`0x109c6ee0`
集体过滤器	`ForEachCollective` @`0x109c7060`(操作码位掩码 `0x1400001340` + `0x56`/`0x5d`/`0x31`)
密钥	`TensorCoreBarrierKey` 演员@`0x109d6200`，`operator<` @`0x109d6620`
着色引擎(同级)	`BarrierColoring<…>::Run` @`0x109cf600` / `0x109d1a60` → 屏障着色
配置运营商	`BackendConfig.BarrierConfig {type @+0x20, id @+0x18, hasbits @+0x10}`
显式环运行时初始化	`ExplicitUniDirRingStrategy::InitializeOnScs` @`0x1337aa60`
策略+0x98回调(普通)	`ExplicitRingRecord` lambda @`0x133a9a40`(`$_2`，种类选择器 `0x4`)
策略+0xb0回调(a2a)	`ExplicitAllToAllRingRecord` lambda @`0x133a94a0`(`$_1`，种类选择器 `0x2`)
回写	`strategy+0x58 = ordinal_`、`+0x60 = next_chip_`、`+0x78 = reordering_map_`
信心	已确认(两个主体均已反编译并重新检查)，除非一行另有说明

1. TC 屏障分配通行证

1.1 `Run` 所在位置 — SC/TC 前叉

BarrierAssignment::RunImpl @0x109c8c00 是共享条目。它接收 XLA 线程池名称和存在线程的分叉的 flat_hash_set<string_view>：

kTensorCoreThread @0x217f78d8 存在 → 从基本配置字段在堆栈上构建 TensorCoreBarrierAssignment 并调用 Run @0x109c8d6e。
kSparseCoreThread @0x217f5ca8 呈现 → 构建一个 SparseCoreBarrierAssignment 并将其命名为 Run @0x109c8e8e。

两个结果具有位或在一起，因此为 TensorCore 线程和 SparseCore 线程编译的模块都运行**两个通道。片上屏障模型是彩色贴图 TC 屏障和静态键控 SC 屏障的结合；两种程序芯片 SFLAG (概述 §1) 的区别仅在于如何选择屏障标识。

text

BarrierAssignment::RunImpl @0x109c8c00
  ├─ if kTensorCoreThread present → TensorCoreBarrierAssignment::Run @0x109c7420   (THIS PAGE §1.2)
  │     coloring ×2 → conflict set → per-key DetermineBarrierConfigForKey → BackendConfig writeback
  └─ if kSparseCoreThread present → SparseCoreBarrierAssignment::Run @0x109c8e8e   (static ring-key dedup)
  result.changed = TC.changed | SC.changed
```text

### 1.2 `TensorCoreBarrierAssignment::Run` @`0x109c7420`

`Run` 是 SparseCore 着色/分配的 TC 对应项。其五个阶段：

1. **构建策略基础。** 构造一个 `AsyncOpPolicy<TensorCoreBarrierKey>` (vtable @`0x217f7a68`)，其中包含兆核字节(`this+0x18`)、核心 ID `std::function<int(HloInstruction*)>` (`this+0x58`)和两个策略布尔值(`this+0x60`、`this+0x68`)。将其重新 vtable 为 `AsyncCollectivePermutePolicy<TensorCoreBarrierKey>` (vtable @`0x217f7ac8`)。
2. **第一次着色通道。** `BarrierColoring<AsyncCollectivePermutePolicy<…>>::Run` @`0x109cf600` → `StatusOr<pair<map<HloInstruction*,long>, flat_hash_set<HloInstruction*>>>`：每操作颜色和冲突操作集。将彩色操作合并到 `std::map<HloInstruction*,long>`(帧 `-0x158`)中，将冲突操作合并到 `FlatHashSet`(帧 `-0x140`)中。
3. **第二次着色过程。** 将策略重新 vtable 到 `AsyncAllToAllWithAsyncBarrierPolicy<TensorCoreBarrierKey>` (vtable @`0x217f7b30`) 并运行 `BarrierColoring<…>::Run` @`0x109d1a60`；将其着色+冲突集合并到**相同** `-0x158` 地图/`-0x140` 集中。
4. **收集集合。** `ForEachCollective` @`0x109c7060` (§1.4) 遍历每个 TC 线程集合，构建每个 `TensorCoreBarrierKey`，并将 `(key, InlinedVector<HloInstruction*, 2>)` 对附加到向量中(条目步长 `0x78`)。 `std::__stable_sort` 与 `Run::$_2` 比较器 @`0x109ccdc0` 使迭代顺序确定。
5. **每键分配循环**(`@0x109c7ed0 … 0x109c8250`、`r12` = 输入、跨步 `0x78`)：
   - 查找合并冲突中条目的第一条指令 `FlatHashSet` (`raw_hash_set::find` @`0x109cf200`) → `has_conflict`;
   - `config = HloModule->config` (`module+0x20`);
   - `DetermineBarrierConfigForKey(&result, key, config, has_conflict)` @`0x109c6fa0` (§1.3) → `BarrierConfig`；
   - **ragged-all-to-all 特殊情况** @`0x109c7f89`：如果操作码 `== ragged-all-to-all(0x56)` 和结果类型 `!= 3` (不是新的每键 id)并且设置了策略 bool `this+0x68` → 转移到 a2a 异步屏障特殊处理(跳转`0x109c8371`);
   - 对于条目 `InlinedVector` 中的 **每个** `HloInstruction`(每个彩色操作共享此密钥)：打开其 `BackendConfig` (`BackendConfigWrapper::GetProto` @`0x1e60dc60`)，设置其 `BarrierConfig` 子消息 = `DetermineBarrierConfigForKey`结果(`BarrierConfig::CopyFrom` @`0x1d6f1fc0`)，设置`BackendConfig` hasbit(`or BYTE [-0x600], 0x10`)，并通过`CloneBackendConfigProto` @`0x1e60dac0` + `BackendConfigWrapper::operator=`重新序列化@`0x1e60de40`。

> **注意 — 按颜色进行重复数据删除，而不是按密钥哈希进行重复数据删除。** 着色引擎证明操作可以共享屏障(在异步依赖关系图中不冲突)获得相同的颜色并且可以共享屏障 id；与 *相同* `TensorCoreBarrierKey` 的操作，但图形冲突被分开。这是与 SparseCore `FlatHashMap` 重复数据删除的结构差异，后者纯粹依赖于静态环配置，并且没有计划重叠的概念。干涉图构建和首次拟合颜色搜索属于[屏障着色](barrier-coloring.md)；该页面仅消耗冲突集为`has_conflict`。

### 1.3 `DetermineBarrierConfigForKey` @`0x109c6fa0` — 种类选择器

有效签名：`BarrierConfig DetermineBarrierConfigForKey(const TensorCoreBarrierKey& key, const HloModuleConfig& config, bool force_global)`。反编译后的主体很短，并且固定了每个商店：

```c
// 0x109c6fa0  (this = &result BarrierConfig; a2 = key; a3 = config; a5 = has_conflict/force_global)
BarrierConfig::BarrierConfig(this, 0);                       // zero-init result
beneficial = IsGlobalBarrierBeneficial(key, config);         // 0x109c6ee0 (§1.5)

if (force_global || beneficial) {                            // → GLOBAL
LABEL_global:
    *(int  *)(this + 0x20) = 1;     // type  = GLOBAL(1)
    *(long *)(this + 0x18) = -1;    // id    = -1 (none)
    *(char *)(this + 0x10) |= 3;    // hasbits: type+id present
    return this;
}

v9 = *(long*)(key + 0x10);          // replica_group_count
if (v9 == *(long*)(key + 0x20) - 1) {           // spans all-but-one group (num_groups - 1)
    if (*(char*)(config + 0x50) != 1) {          // "use global on saturation" flag NOT set
        *(int *)(this + 0x20) = 2;  // type = REPLICA(2)  (shared)
        *(long*)(this + 0x18) = v9; // id   = replica_group_count
        *(char*)(this + 0x10) |= 3;
        return this;
    }
    goto LABEL_global;              // saturation + flag set → GLOBAL(1)
}

// ELSE: provisional fresh per-key
*(int *)(this + 0x20) = 3;          // type = CUSTOM(3)  (fresh)
*(long*)(this + 0x18) = v9;         // id   = replica_group_count (provisional)
*(char*)(this + 0x10) |= 3;
v11 = std::map<key,BarrierConfig>::__try_key_extraction_impl(key, key, config, this);  // 0x109cbd60
if ((dl & 1) == 0)                  // emplace did NOT insert ⇒ key already present (cache HIT)
    BarrierConfig::CopyFrom(this, (BarrierConfig*)(v11 + 0x80));   // adopt the existing config (its type==2)
return this;                        // cache MISS ⇒ keep the fresh type-3 config

数字布局(this+0x20 是字节偏移 0x20 处的 DWORD 类型；this+0x18 是 QWORD id；this+0x10 是 hasbit 字节)：

结果	型 @`+0x20`	id @`+0x18`	时
`GLOBAL`	`1`	`-1`	`force_global` 或 `IsGlobalBarrierBeneficial` 或(饱和度与 `config+0x50 == 1`)
`REPLICA`(共享)	`2`	`replica_group_count`	饱和(`rgc == num_groups − 1`)和`config+0x50 != 1`；或地图重复数据删除发现存在密钥
`CUSTOM`(新鲜)	`3`	`replica_group_count`	否则，在 `std::map` 缓存未命中

GOTCHA — 饱和臂。 饱和情况 (replica_group_count == num_groups − 1) 并不简单地产生 GLOBAL。它产生一个共享的REPLICA(2)(id =真正的replica_group_count)除非模块配置标志config+0x50 == 1，在这种情况下它落入GLOBAL(1)。 0x109c704d 处的 2 写入 (*(int*)(this+0x20) = 2) 是从饱和臂和 std::map 缓存命中 CopyFrom 达到的。 MEGACORE(4) 从未被写入 - 此函数中没有将 movl $4 写入 +0x20(完整 .text 外部引用；概述 §2)。

1.4 `ForEachCollective` @`0x109c7060` — 哪些操作是 TC 集合

ForEachCollective 迭代 HloModule::computations(kTensorCoreThread) @0x10944b40，遍历每条指令，并为操作码字节 (instr+0xc) 位于集合集中的每个操作调用回调。过滤器(@0x109c72b5 … 0x109c736d)：

opcode <= 0x24 上的位掩码 bt 0x1400001340, opcode — 设置位：all-gather(0x06)、all-gather-start(0x08)、all-reduce(0x09)、all-to-all(0x0c)、collective-permute(0x22)、collective-permute-start(0x24)；
cmp 0x56 → ragged-all-to-all 已接受；
cmp 0x5d → reduce-scatter 已接受；
否则(custom-call(0x31))：GetCustomCallCollectiveId @0x109d6540 读取 CustomCallConfig 后端配置字段(hasbit 0x40，返回 +0x78)→ 接受当它携带 TPU 集体 id。

回调还通过递归到操作数位置 (instr+0x40 / instr+0x48)，通过异步开始/异步完成对来线程化冲突图。

字节	助记词	过滤机构	屏障关键作用
`0x06`	全聚	位掩码	副本组密钥
`0x08`	全聚开始	位掩码	副本组键控(异步启动)
`0x09`	全减	位掩码	副本组密钥
`0x0c`	全对全	位掩码	副本组键控；仅操作码符合 `IsGlobalBarrierBeneficial` 条件
`0x22`	集体排列	位掩码	源-目标对键控 (`+0x20`)
`0x24`	集体排列启动	位掩码	源-目标对键控(异步启动)
`0x31`	自定义调用	`GetCustomCallCollectiveId`	接受，当且仅当它具有 TPU 集体 ID
`0x56`	衣衫褴褛	EQ 检查	副本组键控； `+0x51`/`+0x58`特例
`0x5d`	减少散射	EQ 检查	副本组键控(也是 a2a-5d 重写目标)

注意 — 操作码→助记符映射。 该映射是从 HloOpcodeString 长度数组 @0x421c9c0 恢复的(长度 10/16/10/10/18/24/11/17/14 匹配 all-gather / all-gather-start / all-reduce / all-to-all / collective-permute / collective-permute-start / custom-call / ragged-all-to-all / reduce-scatter)加上 .lrodata 字符串存在； char* 表本身已重新定位，因此值通过长度 + 存在来解析。

1.5 `IsGlobalBarrierBeneficial` @`0x109c6ee0` — 启发式

字节精确主体：

// 0x109c6ee0  (key = a1; config = a2)  → bool
if (*(char*)(key + 0x51) == 0)                  // "barrier-heuristic candidate" byte not set
    return false;                               // (VLOG(10) diagnostic path)
if (*(byte*)(key + 0x00) != 0x0c)               // opcode != all-to-all
    return false;
if (*(char*)(key + 0x50) != 0)                  // channel-id parity (cross-module) != 0
    return false;
if (*(long*)(config + 0x170) == 1 ||            // num_partitions / replica_count along axis == 1 (cmpq)
    *(long*)(config + 0x178) == 1)
    return (*(long*)(key + 0x10) == 1);         // replica_group_count == 1 ⇒ single group
return false;
```text

全局屏障对于**单副本组、在拓扑轴上的通道 `all-to-all`** 来说是“有益的”，其测试的分区/副本计数为 `1` — 退化的全对所有，其中每键屏障将同步一个琐碎的组，而单个共享全局屏障更便宜。

> **GOTCHA — `config+0x170` / `config+0x178`.** 两个 `== 1` 测试是 64 位 `cmpq` 比较 (`cmpq $0x1,0x170(%rdx)` / `cmpq $0x1,0x178(%rdx)` @`0x109c6ef3`/`0x109c6efd`)，在这些偏移量处进行字节确认，并在使用时读取为“单设备沿轴”，但它们的原型字段名称(`num_partitions` vs `replica_count` vs `HloModuleConfig` 的派生设备计数)是**推断的**，而不是与描述符进行字段匹配。行为(`==1` ⇒ 测试 `replica_group_count == 1`，本身是 `cmpq $0x1,0x10(%rsi)`)是确定的；字段命名为 LOW。

### 1.6 `TensorCoreBarrierKey` 演员@`0x109d6200` / `operator<` @`0x109d6620`

关键是两个集体必须同意成为共享障碍的“候选者”(然后由颜色决定他们是否真的可以)。 Ctor `TensorCoreBarrierKey(HloInstruction* hlo, function<int(HloInstruction*)> core_id_fn, bool b)`：

| 偏移 | 字段 | 源码 |
|---|---|---|
| `+0x00` | 操作码字节 | `hlo+0xc`;当`all-reduce(0x09)`的计算有≥2个调用者指令时，`all-reduce(0x09)`被重写为`reduce-scatter(0x5d)`(`cmpb $0x9` @`0x109d6297` → `caller_instructions()` ≥ 2 → `movb $0x5d` @`0x109d62db`) |
| `+0x08…+0x18` | 排序 `vector<ReplicaGroup>` | 复制自 `hlo+0xd0`/`hlo+0xd8` (`__assign_with_size` @`0x109d6f60`)，`__introsort`'d @`0x109d7580` (与订单无关)； `+0x10` = 组数 |
| `+0x20…+0x28` | 排序 `vector<pair<long,long>>` | 源-目标对(对于集体排列 `0x22`/`0x24`)：`__assign_with_size` @`0x109d7220` + `__introsort` @`0x109da4e0` |
| `+0x38` | 自定义调用配置标量 | `all-to-all(0xc)` 自定义调用：当 hasbit `0x40` 设置时读取 `CustomCallConfig+0x78`；还设置 `+0x50` true |
| `+0x40` | 核心 ID(整数) | `-1` 默认值，否则为 `core_id_fn(hlo)`(当函数目标非空时通过 `[fn+0x10]` 调用) |
| `+0x50` | 通道 ID 奇偶校验 | `hlo->channel_id() & 1` (`HloInstruction::channel_id` @`0x1e59ff80`) — 字节 `IsGlobalBarrierBeneficial` 门 |
| `+0x51` | 启发式候选标志 | 字节 `IsGlobalBarrierBeneficial` 门控 (`cmpb $0x0,0x51(%rdi)` @`0x109c6ee0`)； **不是由该构造函数编写** - 构造函数仅零接触 `+0x08…+0x37`、`+0x38`、`+0x40`、`+0x48`、`+0x50`、`+0x58`，因此候选字节设置在不同的路径上(站立间隙) |
| `+0x58` | 二次判别式 | `-1` 哨兵仅在 ctor 的 `bool b` 设置且操作码 `== ragged-all-to-all(0x56)` (`test %r15b,%r15b` @`0x109d64dd` → `cmpb $0x56,0xc(%r14)` @`0x109d64e2` → `movq $-1,0x58(%rbx)` @`0x109d64e9`);衣衫褴褛的全局禁用案例 |

`operator<` @`0x109d6620` 按顺序比较：`+0x58`、`+0x38`、`+0x00`(操作码)、`+0x50`(通道奇偶校验)，然后是副本组向量(`+0x10`)计数，然后按组排序的设备 ID)。两个集合共享一个 TC 屏障 **key** iff 相同的操作码、相同的通道奇偶校验、相同的自定义调用标量、相同的辅助判别式以及相同的排序副本组成员身份 - SparseCore 环密钥的密集集合模拟，以 HLO 副本组而不是 ICI 环偏移为键。

---

## 2. `ExplicitUniDirRingStrategy::InitializeOnScs` 及策略+0x98 回调

这是与§1不同的子系统：SparseCore端环策略运行时初始化。当 TC 通道 *决定* 屏障类型时，`InitializeOnScs` *绑定* SC 定序器的屏障 `sync_add` 将针对的每环传输几何结构。它是静态 `IciStrategyRingConfig` EXPLICIT 表 → 编程的 `next_chip` 桥的最后一跳。

### 2.1 `InitializeOnScs` @`0x1337aa60` — 折叠 + 回调调度

System-V 参数：`this` = 策略、`OpBuilder&`、`LocationGenerator`，然后是三个 `Value`(`scs ordinal`、`scs partition`、第三个 `Value` 三元组)。反编译体确认了每一步：

```c
// 0x1337aa60  (a1 = strategy; a2 = OpBuilder; a5 = scs partition Value; a3 = scs ordinal Value)
if (!a5)                                                         // line 351
    return RetCheckFailSlowPath("config != nullptr");

UniDirRingStrategy::InitializeOnScs(a1, a2, a3, a4, ...);        // base init; stores +0x68/+0x70
if (base != OK) return AddSourceLocationImpl(354);

CoreIndex = OffloadFactory::GetCoreIndex(a1+8, a2, a4);          // 0x133e6aa0 (per-replica core index)
v17  = *(a1+8);                                                  // OffloadFactory / Target subobject
v18  = Target::CoresPerChip(v17, /*SC*/2);                       // 0x1d615b40
v19  = Target::LogicalDevicesPerChip(v17, /*SC*/2);             // 0x1d615b00
fold = v18 / v19;                                                // unsigned divide (megacore fold)
v22  = OffloadFactory::IdxConst(a1+8, a2, fold);                 // 0x133e6ba0 (MLIR index constant)
v23  = OffloadFactory::DivU(a1+8, a2, CoreIndex, v22);           // 0x133e6a60 → arith::DivUIOp
ordinal       = v23;
global_core_id = OffloadFactory::ToGlobalCoreId(a1+8, a2, partition, v23);   // 0x133e6880

// THE CALLBACK — call QWORD PTR [strategy + 0x98]
(*(a1 + 0x98))(&record /*StatusOr<ExplicitRingRecord>*/,
               a1 + 0x88 /*captured table descriptor*/,
               a1 + 0x08 /*OffloadFactory subobject*/,
               a2        /*OpBuilder*/,
               &global_core_id,
               &ordinal);

if (record == OK) {                                              // writeback (double-init guarded)
    CHECK(strategy+0x58 == null, "ordinal_ == nullptr");         // line 250
    *(a1 + 0x58) = record.ordinal;                               // ordinal_
    CHECK(strategy+0x60 == null, "next_chip_ == nullptr");       // line 245
    *(a1 + 0x60) = record.next_chip;                             // next_chip_
    CHECK(strategy+0x78 == null, "reordering_map_ == nullptr");  // line 255
    *(a1 + 0x78) = record.reorder;                               // reordering_map_
    return OK;
}
return AddSourceLocationImpl(362);                               // line 362

这三个写回字段在反编译携带的 LogMessageFatal 字符串中命名：ordinal_ at strategy+0x58、next_chip_ at +0x60、reordering_map_ at +0x78 — 每个都受到防范第二次初始化(重新初始化会致命)。 next_chip_ 然后被 next_chip() @0x1337a200 → GlobalCoreIdToPhysicalChipId → ComputeRemoteCoreIndex → SC 自定义屏障发出的 SyncAddOp (屏障 → SFLAG 绑定) 消耗。

注意 — 行号。 RetCheck / AddSourceLocationImpl 行常量(351、354、362)和写回 CHECK 行(245/250/255) 完全匹配 platforms/xla/sparse_core/offload_collective_strategies.cc (立即数 0x15f=351, 0x162=354, 0x16a=362)。 CoresPerChip(2) / LogicalDevicesPerChip(2) 折叠与 ToPartnerGlobalCoreId 使用的超级核心折叠除数相同。

2.2 回调体 — `ExplicitRingRecord` @`0x133a9a40` (plain) / `ExplicitAllToAllRingRecord` @`0x133a94a0` (a2a)

strategy+0x98 插槽由 CreateRingStrategiesForNdFromExplicitTable (emitter_helpers) 内的每颜色工厂封闭件安装。工厂 → 策略图(分解后的 $_N lambda 确认记录类型)：

工厂	构建	回调槽	记录类型	种类选择器
`$_0` @`0x133a9080`	`D2DUniDirRingStrategy`(`0x98` 的尺寸)	—	—	—
`$_1` @`0x133a91e0`	`ExplicitUniDirAllToAllRingStrategy`(`0xe8` 的大小，vtable @`0x21908ec8`)	`+0xb0` = `0x133a94a0`	`ExplicitAllToAllRingRecord`	`0x2`
`$_2` @`0x133a9840`	普通 `ExplicitUniDirRingStrategy`(`0xc0` 的大小，vtable @`0x21908e30`)	`+0x98` = `0x133a9a40`	`ExplicitRingRecord`	`0x4`

注意 - 工厂映射。 分解符号固定三个 lambda 槽：$_0 构建 D2DUniDirRingStrategy，$_1 构建 all-to-all 策略(返回 ExplicitAllToAllRingRecord)，$_2 构建普通显式环(返回ExplicitRingRecord)。 0x133a99d9 上的 $_2 安装将 0x133a9a40 lambda 写入 strategy+0x98 — §2.1 的确切调用目标。

简单查找 ExplicitRingRecord @0x133a9a40 正文：

ctx = [strategy+0x88](表描述符)→ r14：r14[0] = 内部 const-literals 访问器对象(其 operator() 位于 accessor+0x10)； r14[0x10] = id_info_offset (movsxd int32); r14[0x18] = 表基点 ptr； r14[0x20]/r14[0x21] = 字节选择器 + 变量标记。
内部调用：push (kind = 0x4, base, byte-selector, &out, global_core_id, ordinal)； call QWORD PTR [accessor+0x10] @0x133a9a96 — 该访问器对 [base + id_info_offset…] 处的 const-literals int 向量进行切片，并将每核邻居/id/组常量具体化为 MLIR Values。
成功时：next_chip = out[0]、ordinal = out[8 or 0x10](r14+0x21 变体标记选择字段)、reorder = out[0x20] → 写入 StatusOr<ExplicitRingRecord>(record.next_chip +0x8， record.ordinal +0x10、record.reorder +0x18、record.ok +0x0)。释放临时向量。 RetCheck 线 0xfe5。

全部双胞胎 ExplicitAllToAllRingRecord @0x133a94a0 形状相同，但使用种类选择器 0x2 并将结果通过 OffloadFactory::location @0x131e9ca0 + Target::GranuleBytes @0x1d617f80 进行线程化+ BufferOffset::Create @0x133ea5c0 + BufferOffset::WithOffsetElements @0x133eace0 将组信息表切片转换为字节 BufferOffset(每核组信息窗口)。 RetCheck 线 0xfb9。

注意 — 最后一跳。 回调是静态 EXPLICIT 偏移量成为编程 Value 的地方：id_info_offset (table_desc+0x10) + 表基数 (+0x18) 索引运行时常量文字向量，由运行时键入global_core_id / ordinal，用于生成 next_chip 每环 SyncAdd 目标。普通环返回原始 Values； a2a环返回BufferOffset包装的组窗口。 0x4 与 0x2 选择器选择哪个。
INFERRED — 内部访问器主体。 通过 call [accessor+0x10] 到达的访问器对象具有三个经过验证的捕获成员(+0x10 = id_info_offset、+0x18 = 表基数、+0x20/+0x21 = 字节 + 变体)选择器)和固定的调用签名(global_core_id, ordinal, base, selector, kind 0x4/0x2 → out triple)，但它自己的主体 - 执行 spmem 加载的闭包 / DivU-Mod 索引到 const-literals 向量 - 没有单独反汇编。其CALL签名为CONFIRMED； const-literals 向量条目语义是常设生产者间隙 (LOW)。

3.两个子系统，一个芯片SFLAG块

TC 分配通道 (§1) 和显式环 InitializeOnScs (§2) 位于同一片上屏障模型的两端：TC 通道为密集 TC 集合选择 BarrierConfig； InitializeOnScs 绑定 SC 嵌入集合的每环传输几何结构。两者最终都对芯片 SFLAG (屏障 → SFLAG 绑定) 进行编程。

方面	TensorCore (§1)	SparseCore 显式环 (§2)
通过/初始化	`TensorCoreBarrierAssignment::Run` @`0x109c7420`	`ExplicitUniDirRingStrategy::InitializeOnScs` @`0x1337aa60`
标识单元	`TensorCoreBarrierKey`(操作码、通道奇偶校验、副本组/src-tgt 向量、cc 标量)	静态 `IciStrategyRingConfig` 显式偏移量(id_info / group_info)
去重/绑定机制	贪心图着色 (屏障着色) → `BarrierConfig {1,2,3}`	`strategy+0x98` 常量文字查找 → `(next_chip, ordinal, reorder)`
输出目的地	HLO `BackendConfig.BarrierConfig` 子消息	`strategy+0x58/+0x60/+0x78` (`ordinal_` / `next_chip_` / `reordering_map_`)
硬件接收器	芯片SFLAG块通过内核发射器	每环 `SyncAdd` 针对芯片 SFLAG 块上的 `next_chip_`

分割纯粹在于如何选择和绑定屏障身份：TC pass 为实时异步集体冲突图着色；显式环通过运行时核心 ID 索引静态配置表。两个臂均终止于芯片 SFLAG 层 (概述 §1)。

4. 验证说明

libtpu.so v0.0.40 中的字节精确：
DetermineBarrierConfigForKey @0x109c6fa0：BarrierConfig::BarrierConfig(this,0)； IsGlobalBarrierBeneficial； force_global || beneficial→type=1，id=-1；饱和度(key+0x10 == key+0x20 − 1)与config+0x50 != 1→type=2，id=replica_group_count；否则 type=3 新鲜 + std::map __try_key_extraction_impl → 缓存命中 CopyFrom(node+0x80)。类型为 +0x20，id 为 +0x18，拥有 +0x10 |= 3。 没有 movl $4 — 准确。
ExplicitUniDirRingStrategy::InitializeOnScs @0x1337aa60：if (!config) RetCheck 第 351 行；底座 UniDirRingStrategy::InitializeOnScs → AddSourceLocationImpl(354)； GetCoreIndex; CoresPerChip(2) / LogicalDevicesPerChip(2) 无符号除法折叠； IdxConst； DivU； ToGlobalCoreId； call [a1+0x98] 与 (&record, a1+0x88, a1+0x08, OpBuilder, &global_core_id, &ordinal)；写回 *(a1+0x58)=ordinal_、*(a1+0x60)=next_chip_、*(a1+0x78)=reordering_map_，每个 LogMessageFatal 受保护(第 250/245/255 行)； AddSourceLocationImpl(362) — 准确。
符号确认：TensorCoreBarrierAssignment::{Run, ForEachCollective, IsGlobalBarrierBeneficial, DetermineBarrierConfigForKey} 和 TensorCoreBarrierKey 呈现完整的分解符号； ExplicitUniDirRingStrategy::InitializeOnScs 和 CreateRingStrategiesForNdFromExplicitTable $_1 (→ ExplicitAllToAllRingRecord) / $_2 (→ ExplicitRingRecord) lambda 存在，确认工厂→记录映射和 (OffloadFactory&, OpBuilder&, Value, Value) → StatusOr<…Record> 回调签名。
[低]
IsGlobalBarrierBeneficial 的 config+0x170 / config+0x178 标量经过偏移确认并读取为“单设备沿轴”，但原型字段名称(num_partitions 与 replica_count)是推断的，而不是描述符匹配。
内部 const-literals 访问器主体(call [accessor+0x10]，第 2.2 节)- 其三个捕获的成员和调用签名已固定，但对 const-literals 向量进行切片并具体化每核心 Value 的闭包未单独反汇编。
从 __assign_with_size / __introsort 源和操作码分支读取 TensorCoreBarrierKey 子成员语义(其向量是副本组与 src-tgt-pairs)；偏移量已确认，角色命名是结构化的。

交叉引用

屏障算法(本节)

概览 — 基于 SFLAG 的屏障模型、BarrierType 枚举和 InferBarrierConfig 标准化器(本页生产者提要的第二个 BarrierConfig 作者)
屏障着色 — 贪婪干扰图引擎，其合并冲突集是 has_conflict 输入到 DetermineBarrierConfigForKey
屏障 → SFLAG 绑定 — 选定的 BarrierConfig {type, id}(和绑定的 next_chip_)如何成为具体的芯片 SFLAG 编号 + 信号/等待操作
复制品屏障 — 副本组内树屏障 (REPLICA(2))，此传递可以发出的共享 ID 结果

同级子系统

巨核融合 — 巨核折叠集体融合，其环策略 InitializeOnScs 结合
返回索引

二进制： extracted/libtpu-0.0.40-cp314-cp314-manylinux_2_31_x86_64/libtpu/libtpu.so(构建 ID 89edbbe81c5b328a958fe628a9f2207d)。

TensorCore 屏障分配和 InitializeOnScs ​

摘要 ​

1. TC 屏障分配通行证 ​

1.1 Run 所在位置 — SC/TC 前叉 ​

1.4 ForEachCollective @0x109c7060 — 哪些操作是 TC 集合 ​

1.5 IsGlobalBarrierBeneficial @0x109c6ee0 — 启发式 ​

2.2 回调体 — ExplicitRingRecord @0x133a9a40 (plain) / ExplicitAllToAllRingRecord @0x133a94a0 (a2a) ​

3.两个子系统，一个芯片SFLAG块 ​

4. 验证说明 ​

交叉引用 ​

屏障算法(本节) ​

同级子系统 ​