Internal Wiki

主题

InferBarrierConfig ​

地址适用于 libtpu-0.0.40-cp314 轮中的 libtpu.so。其他版本有所不同。

摘要 ​

InferBarrierConfig @0x1376c240 是 钳式融合屏障规范化器：在融合发射时运行的小型决策树，重新读取集体已在其 HLO BackendConfig 中携带的 BarrierConfig，并“降级”每个密钥当 CUSTOM 能够证明集体的通信设备是真正的多参与者时，CUSTOM 就会成为较便宜的 GLOBAL 或 REPLICA 的障碍。它是两个 BarrierConfig 接触点中的第二个——与着色生成器 TensorCoreBarrierAssignment::DetermineBarrierConfigForKey @0x109c6fa0 不同，后者在 HLO 传递时写入原始配置。两者从不共享状态；他们仅通过原型字段进行通信。此页面导出树字节精确并记录树解析的每代 SFLAG 内存映射 — 从芯片配置中雕刻的每个代号基本/计数块 Target::Init / SparseCoreTarget::Init。

该树很短，并且恰好具有三个结构臂。第一个是守卫：针对策略的两轴参与者计数的 单组谓词 a4[1] <= 1 && a4[2] <= 1。单例集不需要重写(仅验证)。第二个臂向其 barrier_type == CUSTOM(3) 的多参与者集合开火：如果集合已通灵，它会咨询 hlo->channel_id() 并使用哨兵 id = -1 重写为 GLOBAL(1)，或者如果集合已通灵，则使用 REPLICA(2)(最后一个可用的 TC 屏障 ID 插槽)重写为 REPLICA(2)。事实并非如此。第三臂是合法性门：BARRIER_INVALID(0) 始终被 RetCheck 拒绝。已设置的 GLOBAL/REPLICA 保持不变； MEGACORE(4) 从未写入。

REPLICA id、Target[+0x8c4] − 1 将树与 SFLAG 块几何结构联系起来：Target+0x8c4 是 TC 屏障 ID 计数，是 Target::Init 从芯片配置中读取的一对 {base, count} 整数之一compiler_reserved 重复的 int32 范围，并且每个保留插槽 SFLAG 公式都是参数化的。此页面记录了每代映射 - 基数/计数雕刻和独立于生成的块结构 - 因此重新实现者可以解析 Target[+0x8c4] − 1 和五个保留的顶部插槽，而无需重新跟踪芯片配置。 BarrierType 枚举、着色引擎、SFLAG 数字绑定和每个代号整数的文字都是同级页面(链接如下)；此页面拥有 决策树和 SFLAG 基数/计数图。

重新实现，合约为：

  • 三臂决策树 - 单组谓词 a4[1] <= 1 && a4[2] <= 1、CUSTOM(3) → 通道 id 门控 GLOBAL/REPLICA 重写，以及 INVALID(0) RetCheck - 具有精确的本地布局(type @-0x30、id @-0x38、hasbits @-0x40)，因此重写写入正确的字节。
  • 两个重写目标： GLOBAL(1)携带id = -1； REPLICA(2) 带有 id = Target[+0x8c4] − 1，顶部可用的 TC 屏障 ID 插槽。
  • 每代 SFLAG 内存映射 树将 Target[+0x8c4] 解析为：{base @+0x8c0, count @+0x8c4} TC 块 (count = |CR_TC| − 5)、{base @+0x1d0, count @+0x1d4} SC 块(count = |CR_SC|，无 −5)以及五个保留的块TC 顶部插槽。
功能xla::jellyfish::(anonymous namespace)::InferBarrierConfig @0x1376c240
签名Status InferBarrierConfig(Target const&, HloInstruction const*, Strategy*) → StatusOr<BarrierConfig>
TUplatforms/xla/service/jellyfish/lowering/rotated_pincer_fusion_emitter.cc
呼叫者8，所有 RotatedPincer / RotatedPincerShort / AsyncPincer 融合发射器
谓词*((int64*)a4 + 1) <= 1 && *((int64*)a4 + 2) <= 1(策略+0x8 / +0x10)
定制臂type==3 → 引导？ GLOBAL(1), id=-1 : REPLICA(2), id=Target[+0x8c4]−1
从不写入MEGACORE(4) — 身体任何地方都没有 movl $4
TC块Target+0x8c0 底座 / Target+0x8c4 计数 (= |CR_TC| − 5)
SC块SparseCoreTarget+0x1d0 底座 / +0x1d4 计数(= |CR_SC|，无 −5)

枚举+生产者→标准化器→降流在障碍 — 剖面图； 屏障着色上的着色引擎； 屏障到 SFLAG 的绑定 上的 SFLAG 编号降低； 专用同步标志 上的 compiler_reserved 原型源； 按代号编译器保留 上的字面每代整数。此页面不会重复它们。

1. 决策树 ​

用途 ​

从一个 HLO 集体发射钳形融合，其屏障在 HLO 通过时已着色。在发射时，融合的“实际”通信集形状是已知的——参与者在其两个轴上进行计数——这是着色通道所没有的。 InferBarrierConfig 利用了这一点：如果集合是真正的多参与者，则每个密钥的 CUSTOM 屏障是浪费的(它为可以共享设备全局或副本组插槽的屏障烧毁了新的 SFLAG id)，因此树将其折叠为更便宜的共享类型。如果该集合是退化的(两个轴上都是单一的)，则显式的 CUSTOM 着色将保持不变。该函数是标准化器，而不是生产者：它只会降级 CUSTOM，并且永远不会触及已经设置的 GLOBAL/REPLICA。

入口点 ​

text
RotatedPincerFusionEmitter / RotatedPincerShortFusionEmitter / AsyncPincerFusionEmitter
  ├─ CreateEmitterForMultipleInputsOrOutputs    @0x1376b96c
  ├─ EmitSingleInputAllReduceScatterFusion      @0x1376dac9
  ├─ EmitSingleInputAllGatherFusion             @0x13773230
  ├─ EmitColorwiseFusedAllReduces (Rotated)     @0x13774e15
  ├─ EmitColorwiseFusedAllReduces (Short)       @0x1377554e
  ├─ EmitSingleInputAllReduceScatterFusion (A)  @0x137760a0
  ├─ EmitAllReduceScatterFusion (Async)         @0x137775d3
  └─ EmitColorwiseFusedAllReduces (Async)       @0x13778359
        └─ InferBarrierConfig(target, hlo, strat)   @0x1376c240   ── this page
```text

所有八个调用者都是通过完整的 `.text` 外部参照扫描解析的直接 `E8 rel32` 调用；没有虚拟调度。 `InferBarrierConfig` 位于 `rotated_pincer_fusion_emitter.cc` 的文件本地匿名命名空间中，因此它没有外部链接 - 它只能从该 TU 中的八个钳形发射器访问。

### 算法

身体较短。下面是带注释的树； IDA 本地变量(`v35` = type、`v34` = id、`v33` = hasbits、`a4` = Strategy、`a2` = Target、`a3` = HLO)被命名，并引用反编译器行号和帧偏移以进行交叉检查。

```c
function InferBarrierConfig(target /*a2*/, hlo /*a3*/, strat /*a4*/):   // 0x1376c240
    cfg = hlo->backend_config<BackendConfig>()                  // 0xf58e6c0; StatusOr @-0x490
    if (!cfg.ok())                                              // line 37
        return RetCheck(line 93, rotated_pincer_fusion_emitter.cc)   // AddSourceLocationImpl, line 41

    BarrierConfig bc;                                           // local @-0x50 (ctor line 71)
    if (cfg.has_barrier_config())                               // hasbit 0x10 in v27 @-0x268, line 72
        bc.CopyFrom( cfg.barrier_config() ?: BarrierConfig_globals_ )  // default @0x223a9450, lines 74-77
    // bc fields now: type = v35 @-0x30, id = v34 @-0x38, hasbits = v33 @-0x40

    // ARM 1 — singleton-group predicate (the GUARD)
    if (strat[1] <= 1 && strat[2] <= 1):                        // *((int64*)a4+1)<=1 && *((int64*)a4+2)<=1, line 79
        type = bc.type                                          // v14 = v35, validate-only (line 81)
        // fall through to legality gate (ARM 3)
    else:                                                       // MULTI-participant on either axis
        type = bc.type                                          // v14 = v35, line 85
        // ARM 2 — CUSTOM downgrade
        if (bc.type == CUSTOM /*3*/):                           // if (v35 == 3), line 86
            hlo->channel_id()                                   // 0x1e59ff80; has_value → v15 (dl), line 88
            if (v15 == 1):                                      // channelled collective, line 89
                bc.type = GLOBAL  /*1*/;  bc.id = -1            // v35=1, v16=-1, lines 91-92
            else:                                               // non-channelled
                bc.type = REPLICA /*2*/                         // v35=2, line 97
                bc.id   = target[0x8c4] - 1                     // v16 = *((int*)a2 + 561) - 1, line 99
            bc.hasbits |= 3                                     // v33 = v17 | 3, line 102
            goto WRITE                                          // LABEL_24, line 103
        // else (already GLOBAL/REPLICA, type != 3) → keep; fall to ARM 3

    // ARM 3 — legality gate (reached by singleton OR multi-but-not-CUSTOM)
    if (type != 0):  goto WRITE                                 // if (v14) goto LABEL_24, line 146
    return RetCheck(line 115,                                   // BARRIER_INVALID rejected, line 148
                    "barrier.barrier_type() != BarrierType::BARRIER_INVALID")

WRITE:                                                          // LABEL_24, line 103
    out.barrier_config = bc                                     // BarrierConfig ctor+CopyFrom into this+8
    *(int64*)this = 1                                           // StatusOr "ok" tag, line 109
    return OK

注意 — RetCheck 调用中的源行号是十进制 93 和 115(原始立即数中的 0x5d 和 0x73)，两者都引用相同的两个 rotated_pincer_fusion_emitter.cc 行。 INVALID RetCheck 绳子 "barrier.barrier_type() != BarrierType::BARRIER_INVALID" 是唯一按名称固定 BARRIER_INVALID = 0 的 .rodata 锚。

三个臂，重述 ​

该树恰好具有三个可供重新实现者重现的结构结果：

臂进入条件行动
1 — 单例防护strat[1] <= 1 && strat[2] <= 1(双轴≤1)没有重写。降至 ARM 3(仅验证)。
2 — 自定义降级多参与者 和 type == CUSTOM(3)通道 → GLOBAL(1), id=-1；否则REPLICA(2), id=count−1； hasbits |= 3；写。
3 — 合法性门通过单例或通过多个但已经 GLOBAL/REPLICA 达到type != 0 → 保存并写入； type == 0 (INVALID) → 重新检查。

QUIRK — CUSTOM(3) 测试仅位于多参与者 else 分支内部。 单例通信集上的 CUSTOM 屏障永远不会降级 - 它属于 ARM 3，其中 type == 3 不为零，因此它会逐字保留。 CUSTOM 成为 GLOBAL/REPLICA 的唯一方式是严格的多参与者集。将 type == 3 检查提升到谓词之上的重新实现会错误地将退化的单核集合体折叠到共享屏障上。

明白了 — 重写是对枚举值进行的单向降级，而不是一般的重新映射。 type == 3 是触发新类型写入的唯一值； type == 1/type == 2 直接通过 ARM 3 不变。 没有 movl $4，也没有产生 MEGACORE(4) 的路径——通过阅读全文确认。关闭所有五个 BarrierType 值的开关在此处具有死区 case 4。

2.谓词——策略的两轴参与者计数 ​

它读什么 ​

防护是两个独立的 int64 与常数 1 进行比较：

c
if ( *((__int64 *)a4 + 1) <= 1 && *((__int64 *)a4 + 2) <= 1 )   // line 79
```text

`a4` 是 `Strategy*`。这两个字段是 `Strategy+0x8` (`a4[1]`) 和 `Strategy+0x10` (`a4[2]`)。与 `cmpq $1` 相比，两者都是 64 位有符号负载 — `a4[1] > 1` 和 `a4[2] > 1` 是多参与者条件，ORed(反编译器将 De Morgan 对偶呈现为 ≤-1 测试的 `&&`)。 **两个**轴上均为 `1`(或 `0`)的集合是“单例”； **任一**轴上的 `> 1` 是“多参与者”。

### 轴的含义

这两个字段是钳形集体在其两个轴上的通信集的参与者计数 - StrategyND 发射器在调用发射器之前从 HLO 的 `replica_groups` / `ShardingConfig` 构建的第 0 阶段/第 1 阶段环或副本组大小。 “多参与者”(任一轴上的 `> 1`)意味着需要真正的跨核心交会，因此每个键的 `CUSTOM` 着色被折叠为更便宜的共享屏障； “单例”意味着集体是退化/单核的，并且保留显式的彩色屏障。

> **GOTCHA —** 偏移量 `Strategy+0x8` / `Strategy+0x10` 是字节确认的(每个都是 `cmpq $1`)，但字段 *names* (哪个轴是第 0 阶段环长度与第 1 阶段副本组计数)来自 StrategyND 融合上下文，**不是**来自结构描述符。将“环长度与副本组计数”读数视为推断值(中)。 *行为* - 任一轴上的 `> 1` 都会触发降级 - 是确定的。设置 `+0x8`/`+0x10` 的策略编写者(例如 `StrategyND::BuildStrategy` 和 `GetPhase{0,1}ReplicaGroups` 帮助程序)未针对此页面进行单独跟踪。

---

## 3. 两个重写目标

当 ARM 2 触发时，树将两个 `{type, id}` 对之一写入本地 `BarrierConfig`(`type @-0x30`、`id @-0x38`、`hasbits @-0x40`)，然后写入 `hasbits |= 3` 来标记两个字段都存在。

### GLOBAL — 通道集体

```c
xla::HloInstruction::channel_id(a3);   // 0x1e59ff80
if (v15 /*has_value, dl*/ == 1) {
    v35 = 1;     // BarrierType::GLOBAL @-0x30
    v16 = -1;    // id = -1 (the sentinel GLOBAL id) @-0x38
}

集合上出现的 channel_id 标志着它是一个跨模块/跨设备 通道主导 集合，它必须在设备范围的全局屏障上会合。该树写为GLOBAL(1)，哨兵为id = -1；降低从保留的顶部插槽 base+count+4(GetGlobalBarrierSyncFlagNumber，§4)解析 GLOBAL 屏障的 SFLAG 编号，忽略 -1 id 占位符。 channel_id @0x1e59ff80 返回 optional<int64> 与 value = *(hlo+0xc0)，has_value = *(hlo+0xc8) 加载到 dl 中并与 == 1 进行比较。

REPLICA — 非通道集体 ​

c
else {
    v35 = 2;                          // BarrierType::REPLICA @-0x30
    v16 = *((int *)a2 + 561) - 1LL;   // id = Target[+0x8c4] - 1 @-0x38
}
```text

否 `channel_id` 表示副本组内集合体，仅需要副本组树屏障。该树写入 `REPLICA(2)` 并将 id 固定到 `*((int*)a2 + 561) − 1` = `Target+0x8c4 − 1` = `count − 1` — 保留块中**最后一个可用的 TC 屏障 ID 槽**(`561 * 4 = 0x8c4`；请参阅§4 了解为什么 `Target+0x8c4` 是计数)。这是决策树依赖于每代 SFLAG 映射的单点：`REPLICA` id 不是新分配，而是对可用 id 窗口顶部的固定引用。

> **QUIRK —** `REPLICA` id 固定到 `count − 1`，每个钳形融合的 `REPLICA` 降级都位于同一插槽上。这是有意的共享：来自该标准化器的所有副本组钳形屏障都重复使用一个 id(可用窗口的顶部)，与着色生产者的 `REPLICA` id(*每个键*共享)不同。重新实现者不得在此处分配新的 id — 该值是芯片配置计数的固定函数。

---

## 4. 每代 SFLAG 内存映射

决策树的 `REPLICA` id (`Target[+0x8c4] − 1`) 和每种屏障类型的降低都针对在目标 init 处雕刻的每核类型**保留的 SFLAG 块**进行解析。本节记录了树所依赖的块几何形状。区块(`compiler_reserved`)的原型*来源*位于[专用同步标志](special-purpose-sync-flags.md)； id 与 SFLAG memref 的*数字绑定*位于 [屏障到 SFLAG 的绑定](barrier-to-sflag-binding.md) 上；每个代号整数的文字位于 [按代号编译器保留](per-codename-compiler-reserved.md) 上。此页面拥有**基础/计数雕刻和保留槽位地图**。

### 雕刻

`Target::Init` @`0x1d60fc20` 从芯片配置中复制 `compiler_reserved(TensorCore)` 重复的 int32 范围，检查它是连续升序的，然后写入：

```c
*((_DWORD *)target + 560) = arr[0];        // base  → Target+0x8c0   (560*4 = 0x8c0)
*((_DWORD *)target + 561) = size - 5;      // count → Target+0x8c4   (561*4 = 0x8c4)

count = size − 5 为指定的跨核心屏障同步标志保留范围的前 5 个(如下)。 SparseCoreTarget::Init @0x1d612b20 对 compiler_reserved(SparseCore) 执行相同的操作，但没有 −5：

c
*(_DWORD *)(sctgt + 464) = SpecialPurposeSyncFlags[1];      // SC base  → SparseCoreTarget+0x1d0
*(_DWORD *)(sctgt + 468) = SpecialPurposeSyncFlags->size;   // SC count → SparseCoreTarget+0x1d4 (FULL, no −5)
```text

SC 块已满，并在 *`[SC_base, SC_base+SC_count)` 内保留其全局屏障 id。 TC 和 SC 范围在构造上是不相交的 - 它们是不同的 `SpecialPurposeSyncFlags` 原始消息，由不同的 `TpuCoreType` 键入，从 `GetSpecialPurposeSyncFlags(core)` @`0x20afcf40` 读取(索引 `core << 6`，即 `+0x2a0 + core*0x40`；TensorCore 条目是强制性的或`Target::Init` 通过 `DieBecauseNull` 消亡)。

| 块 | 基场 | 计数字段 | 计数公式 |
|---|---|---|---|
| TensorCore | `Target+0x8c0` (`target[560]`) | `Target+0x8c4` (`target[561]`) | `\|CR_TC\| − 5` |
| SparseCore | `SparseCoreTarget+0x1d0` (`+464`) | `SparseCoreTarget+0x1d4` (`+468`) | `\|CR_SC\|` (no `−5`) |

> **明白了 —** `SparseCoreTarget+0x90` **不是** SFLAG 窗口基础：`*(sctgt+144) = TpuCoreParts::SequencerCount(core, 5)`，每个核心定序器计数，而不是屏障 id。 `SparseCoreTarget+0x1fc` (`*(sctgt+508) = v77 − 4`) 是 `GetMemoryReservation → GetUserRegion` 长度(Mosaic 每核树屏障窗口，MemorySpace 14)，第三个不相交区域。两者都不是 `compiler_reserved` SFLAG 块的一部分。 SC 树屏障窗口位于 [树屏障/vSync](tree-barrier-vsync.md)。

### 预留的 5 个 TC 顶部插槽(`−5`)

`count = size − 5` 将 TC 范围的前 5 个 id 雕刻到指定的跨核心屏障同步标志中。所有三个访问器公式都是字节精确的(`target[560]` = `Target+0x8c0` = 基数；`target[561]` = `Target+0x8c4` = 计数)：

| 插槽 | 配件 | 公式 |
|---|---|---|
| `base + count + 0` | `GetMegacoreBarrierSyncFlagNumber` @`0x1d60f4e0` | `target[560] + target[561]` = `base + count`(`Megacore()` 门控；检查 `"topology_->chip_config().Megacore()"`，第 154 行) |
| `base + count + 1` | (间隙；`GetAllReduceSyncFlagNumber(0)` 非法 — 检查 `phase > 0`) | `base + count + 1`(永久间隙) |
| `base + count + 2` | `GetAllReduceSyncFlagNumber(1)` @`0x1d60f440` | `target[560] + 1 + target[561] + 1` = `base + count + 2` |
| `base + count + 3` | `GetAllReduceSyncFlagNumber(2)` @`0x1d60f440` | `base + count + 3` |
| `base + count + 4` | `GetGlobalBarrierSyncFlagNumber` @`0x1d60f420` | `target[561] + target[560] + 4` = `base + count + 4` |

`GetAllReduceSyncFlagNumber(phase)` 是 `LogMessageFatal` 绑定到 `0 < phase < 3` (检查第 143/144 行)，这就是为什么 `base+count+1` 是永久间隙 - `phase = 0` 是非法的，因此调用者无法命名它。 **可用的每个 id 窗口**是 `[base, base+count)`：`REPLICA` 和 `CUSTOM` id 满足 `id < count`，严格位于五个保留插槽的“下方”。这正是决策树的 `REPLICA` id、`count − 1`、索引的窗口 - 可用范围的顶部，第一个保留槽的下方。

```text
TC SFLAG block (Target+0x8c0 = base, Target+0x8c4 = count):

  base                          base+count                      base+count+5
   |  usable per-id window [base, base+count)  |   5 reserved top slots   |
   |  CUSTOM / REPLICA ids (id < count)         |  mega gap ar1 ar2 glob   |
   |                          ^                 | +0   +1  +2  +3  +4      |
   |                          |                 |
   |          REPLICA id = count-1 (normaliser, §3) lands here

每代结构(参数) ​

让 CR_TC = compiler_reserved(TensorCore) 和 CR_SC = compiler_reserved(SparseCore) 用于给定的 (codename, deployment-name) 芯片配置。每一代的块结构都是相同的——只有整数不同，因为 −5 是 Target::Init 中的编译时常量(反编译器将其渲染为 size − 5；反汇编为 add $0xfffffffb)，并且每个保留槽公式在 (base, count) 中都是参数化的。

Gen(代号)TC块(+0x8c0/+0x8c4)SC块(+0x1d0/+0x1d4)TC top-5 保留(块内)
JF(kJellyfish，v2)base=CR_TC[0]，count=|CR_TC|−5base=CR_SC[0]，count=|CR_SC|巨型 b+c，间隙 b+c+1，ar1 b+c+2，ar2 b+c+3，glob b+c+4
DF(kDragonfish，v3)相同相同同5槽图
PF(kPufferfish，v4)相同相同同5槽图
VF(kViperfish，v5p)相同相同同5槽图
GL(kGhostlite，v6e)相同相同同5槽图
GF(k6acc60406，v7)相同相同同5槽图

Megacore 部署(megacore*、megachip)是指 CoresPerChip(TensorCore) == 2 → BarrierMegacore 处于活动状态并且消耗 base+count megacore 插槽的部署；其他部署将其保留为未使用(GetMegacoreBarrierSyncFlagNumber 上的 Megacore() 门失败)。每代表是一个结构，而不是一个值表 - 文字 CR_TC[0] / |CR_TC| / CR_SC[0] / |CR_SC| 整数是嵌入式内存文件依赖项。

明白了 — 每个 (codename, deployment-name) 整数的文字 不能从 .rodata 静态提取。它们位于嵌入式芯片配置 memfile 二进制 pb blob (tpu_chip_config_memfile_{default,megacore,megachip,…}_embed_internal_create @0x20b18fa0..) 中，在运行时通过 FLAGS_deepsea_chip_config_name @0x224714b0 键控的 flat_hash_map<tuple<TpuVersion, name, TpuCoreType>, FileToc*> 进行解析。区块几何形状和−5已确认；整数为低(memfile 依赖性)。请参阅 按代号编译器保留。

5. 与着色生产商的关系 ​

InferBarrierConfig 是 BackendConfig.BarrierConfig 领域的两位作者之一；重新实现者必须使它们保持独特。

制片人运行时写入id 来源
DetermineBarrierConfigForKey @0x109c6fa0HLO 屏障分配通行证(每个键)GLOBAL(1) / CUSTOM(3 fresh) / REPLICA(2 shared)-1(全球)/新鲜/共享密钥 ID
InferBarrierConfig @0x1376c240钳式融合发射(每次融合，8 个调用者)CUSTOM → GLOBAL(1, id=-1)(如果有通道)； CUSTOM → REPLICA(2, id=count−1) 如果没有-1(全球)/ count−1(复制品)

DetermineBarrierConfigForKey 是 HLO 通行时的权威着色生成器：它运行每键冲突/着色图(由 屏障着色 提供)并分配原始 {type, id}。 InferBarrierConfig 是一个灌注 标准化器，仅在发射时触发，并且仅降级 CUSTOM 选择 - 它从不升级，从不重写已设置的 GLOBAL/REPLICA，并且从不写入 MEGACORE(4)。两个结果都提供相同的降低 (屏障到 SFLAG 的绑定) → 相同的每代 SFLAG 数字空间 (§4)：GLOBAL 解析为 base+count+4，REPLICA 解析为 base+id(来自此标准化器的 id = count−1，顶部可用插槽)。

注意 — 分工是“颜色，然后专业”。着色通道在不知道每个融合的参与者集的情况下决定整个模块的障碍“共享”；一旦钳融合的实际环/复制品组形状实现，标准化器就会专门化幸存的 CUSTOM 选择。两者都不是完整的障碍分配故事 - 当涉及钳形融合时，内核最终降低的 BarrierConfig 是标准化器的输出，否则是着色的输出。

六、验证注意事项 ​

libtpu.so v0.0.40 中的字节精确：

  • InferBarrierConfig @0x1376c240 全文：谓词 *((__int64*)a4+1) <= 1 && *((__int64*)a4+2) <= 1(第 79 行)； if (v35 == 3) 定制(第 86 行)； channel_id(a3) 然后 v15 == 1 → v35=1, v16=-1 GLOBAL (第 88-92 行) else v35=2, v16 = *((int*)a2+561) − 1 REPLICA (第 97-99 行)； v33 = v17 | 3 hasbits(第102行)； if (v14) keep-else-RetCheck 第 115 行 "barrier.barrier_type() != BarrierType::BARRIER_INVALID"(第 148 行)；默认BarrierConfig_globals_ @0x223a9450(第76行)； 没有 movl $4 — 准确。
  • Target::Init @0x1d60fc20：target[560] = arr[0](基础→ +0x8c0)，target[561] = size − 5(计数→ +0x8c4) - 精确(反编译行2067-2068)。
  • SparseCoreTarget::Init @0x1d612b20：(sctgt+464) = SPSF[1](SC 底座 → +0x1d0)、(sctgt+468) = SPSF->size(SC 计数 → +0x1d4，无 −5)； (sctgt+144) = SequencerCount(core, 5)(+0x90，不是 SFLAG)— 准确。
  • GetGlobalBarrierSyncFlagNumber @0x1d60f420：target[561] + target[560] + 4 = base + count + 4 — 准确。
  • GetAllReduceSyncFlagNumber @0x1d60f440：检查 phase > 0 (143) / phase < 3 (144)； target[560] + phase + target[561] + 1 — 准确。
  • GetMegacoreBarrierSyncFlagNumber @0x1d60f4e0：Megacore() 门控； target[560] + target[561] = base + count — 准确。
  • GetSpecialPurposeSyncFlags @0x20afcf40：_bittest64(*(chip+864), core) 门； core >= 3 → ud1； return chip + 672 + (core << 6) = +0x2a0 + core*0x40 — 准确。
  • 8 个呼叫者，所有钳形融合发射器，由 E8 rel32 外部参照扫描 — 精确。

[中] 策略 +0x8 / +0x10 字段名称(哪个轴是第 0 阶段环长度与第 1 阶段副本组计数)来自 StrategyND 融合上下文；偏移量和 cmpq $1 谓词是确定的，名称是推断出来的。

[低] 每代 compiler_reserved 整数的文字 (§4) — proto 字段、carve 公式和 memfile 查找已确认，但整数是从嵌入的 binarypb blob 运行时解析的，并且不是静态提取的。 BarrierType 数字 2 (REPLICA) 和 4 (MEGACORE) 从 movl/cmp 字节模式恢复；仅 INVALID(0)、GLOBAL(1)、CUSTOM(3) 显示为命名的 .rodata 字符串。

交叉引用 ​