巨核集体融合

此页面上的所有地址、符号和偏移量均适用于 libtpu-0.0.40-cp314 轮中的 libtpu.so(构建 ID 89edbbe81c5b328a958fe628a9f2207d，781,691,048 字节，未剥离，.text VMA == 文件偏移量，构建libtpu_lts_20260413_b_RC00)。其他版本会有所不同；将每个 VA 视为版本固定。

摘要

megacore 芯片带有每个逻辑设备两个 TensorCore，编译器将其折叠到单个 XLA 设备中。因此，由该逻辑设备发出的集体必须被分割到两个物理核心(偶/奇对)，然后重新同步，以便在结果离开芯片之前两半一致。此页面拥有重新实现者无法从 op 表面或每轴环计划恢复的折叠的三个部分：

MEGACORE(4) 标记从未实现。 BarrierType 有五个臂，但只有 GLOBAL(1) / REPLICA(2) / CUSTOM(3) 被写入 BarrierConfig； BARRIER_INVALID(0) 是全零默认值。 MEGACORE(4) 从未生成 - 没有编译器传递写入 barrier_type == 4，并且 net_util::GetBarrierSyncFlag (0x1c69ad00) 的 type-4 臂在结构上已死亡。实际的 TensorCore 巨核屏障是一个 运行时 构造 - LloRegionBuilder::BarrierMegacore (0x1d5222a0) - 由 TpuChipConfig::Megacore() 门控，完全在 BarrierConfig → GetBarrierSyncFlag 路径之外。
两核折叠。 BarrierMegacore 及其程序描述符双胞胎 net_util::SynchronizeProgramDescriptorStatesMegacore (0x1c697540) 各自构建一个 Fingerprint2011 派生的屏障 id，然后在保留超级核心上运行 CoreIndex → SmodU32(core_count=2) → SaddS32 偶/奇同步添加SFLAG 插槽(Target::GetMegacoreBarrierSyncFlagNumber 0x1d60f4e0，地址标签 "megacore barrier sync flag")，CHECK(core_count == 2) / CHECK(TensorCoresPerLogicalDevice() == 2) 强制偶/奇对。
TC SFLAG 块初始化。 Target::Init (0x1d60fc20) 从芯片的专用 SFLAG 范围中雕刻出 megacore / AllReduce / Global 访问器：base = arr[0] → Target+0x8c0、count = arr.size() − 5 → Target+0x8c4。 −5 为指定的运行时访问器保留前五个插槽，其中 base+count+0 是超级核心插槽。

偶/奇扭曲拓扑(哪个物理核心是“主要”还是“辅助”，两个核心如何映射到圆环上)位于 Megacore 偶数/奇数分割； SFLAG 握手协议(VsyncAdd / VwaitGeSV 语义，完整的 BarrierType 分类法)位于屏障 / TensorCore 屏障 / SpecialPurposeSyncFlags 中。本页重点介绍*MEGACORE(4) 标记、两核集体折叠和超级核心SFLAG 插槽初始化*。


Megacore TC 屏障(运行时折叠)	`LloRegionBuilder::BarrierMegacore` — `0x1d5222a0`
程序描述符状态折叠	`net_util::SynchronizeProgramDescriptorStatesMegacore` — `0x1c697540`
Megacore SFLAG 插槽配件	`Target::GetMegacoreBarrierSyncFlagNumber` — `0x1d60f4e0` (`base+count`)
SFLAG-块填充	`Target::Init` — `0x1d60fc20`(`+0x8c0 = arr[0]`、`+0x8c4 = size − 5`)
专用 SFLAG 源	`TpuChipConfig::GetSpecialPurposeSyncFlags(TpuCoreType)` — `0x20afcf40`
死亡型4臂	`net_util::GetBarrierSyncFlag` — `0x1c69ad00`
百万核模式门	`TpuChipConfig::Megacore()`; `TpuChipConfig+0x7c ≥ 2`； `CoresPerChip(TC) == 2`
来源 TU	`platforms/xla/service/jellyfish/target.cc`、`.../llo_region_builder.cc`、`.../lowering/net_util.cc`

1. `MEGACORE(4)` 标记绝不是 `BarrierConfig`

BarrierType 是一个 5 值枚举，在结构偏移 +0x20(字段 1)处存储为 BarrierConfig.barrier_type。这种折叠的决定性结构结果：五个臂中只有三个被序列化为 BarrierConfig。

`BarrierType`	作为 `BarrierConfig` 制作？	实现为	SFLAG 插槽
`BARRIER_INVALID(0)`	默认值(全零，从未写入)	`GetBarrierSyncFlag` 检查失败目标	—(非法)
`GLOBAL(1)`	是 — `DetermineBarrierConfigForKey` / `InferBarrierConfig`	`GetGlobalBarrierSyncFlagNumber`	`base+count+4` (TC)
`REPLICA(2)`	是 — 相同的生产商	每个 id 窗口	`base + id` (`id < count`)
`CUSTOM(3)`	是 — 相同的生产商	每个 id 窗口	`base + id` (`id < count`)
`MEGACORE(4)`	NO — 从未在任何地方写过	运行时`BarrierMegacore`(不是配置)	`base+count`(单独的附件)

写入 BarrierConfig.barrier_type 的两个函数 — DetermineBarrierConfigForKey(0x109c6fa0，仅写入 1 / 2 / 3)和标准化器 InferBarrierConfig (0x1376c240，在单例通道上重写 CUSTOM→GLOBAL，可以设置 REPLICA，从不设置 4)——两者都被确认永远不会发出 4。 没有特定于巨核的分配通道。 MEGACORE 枚举器是一个永远不会具体化到配置中的枚举值。

1.1 `GetBarrierSyncFlag`的死4型臂

net_util::GetBarrierSyncFlag (0x1c69ad00) 是将 BarrierConfig 降低到 SFLAG 编号。反编译显示barrier_type(a1+0x20，即*(int*)(config+32))恰好驱动两个分支加上一个CHECK - 2 vs 3 vs 4没有每种类型的分割：

// net_util::GetBarrierSyncFlag — 0x1c69ad00 (decompiled, condensed)
v3 = *(_DWORD *)(a1 + 32);             // barrier_type (field 1, offset 0x20)
if (v3 == 1) {                          // GLOBAL
    n = Target::GetGlobalBarrierSyncFlagNumber(target);   // base+count+4
    return SflagImmPtr(n, "global barrier sync flag", 24);
}
if (!v3)                                // BARRIER_INVALID
    CHECK("barrier.barrier_type() != BarrierType::BARRIER_INVALID");  // net_util.cc:2065
// ---- ALL OF {2,3,4} share this one arm: ----
v4 = *(_QWORD *)(a1 + 24);              // barrier.id()
CHECK(v4 < *(int*)(target + 2244));     // id < GetBarrierSyncFlagCount()  (target+0x8c4)  net_util.cc:2070
n = v4 + *(_DWORD *)(target + 2240);    // base(+0x8c0) + id
return SflagImmPtr(n, "barrier sync flag number", 24);
```text

`2240 = 0x8c0`(基本)、`2244 = 0x8c4`(计数)。 `REPLICA(2)`、`CUSTOM(3)` 和无法访问的 `MEGACORE(4)` 将**全部**解析为与 `base + id`(每个 id 窗口)、绑定检查的 `id < count` 相同。没有字节来区分它们——每种类型的语义完全位于*哪个发射器运行*的上游。由于两个制作人都没有编写过 `4`，所以这个手臂对于 type 4 来说是死的，并且 - 至关重要的是 - 即使*如果*合成了 type-4 配置，它也会解析为 `base + id`，**不是**保留的 `base+count` 超级核心插槽。正是出于这个原因，巨核屏障故意绕过 `BarrierConfig`。

> **注意 —** 完整的每个代号 `BarrierConfig` 生产者/消费者分类(`DetermineBarrierConfigForKey` 着色通道、`InferBarrierConfig` 标准化器、GLOBAL `base+count+4` 插槽、每个 id 窗口)位于 [屏障](../barrier/overview.md) 中， [InferBarrierConfig](../barrier/infer-barrier-config.md) 和 [SpecialPurposeSyncFlags](../barrier/special-purpose-sync-flags.md)。此页面只需要结果：类型 4 不是配置。

---

## 2. TensorCore Megacore 屏障 (`BarrierMegacore`)

实际的 TensorCore megacore 屏障由 LLO 区域生成器直接发出。 `LloRegionBuilder::BarrierMegacore`(`0x1d5222a0`)将一个逻辑设备的两个物理TensorCore折叠在预留的megacore SFLAG插槽上。

### 2.1 巨核模式门

该函数是无操作的，除非超级核心条件的**所有三个**都成立(反编译)：

```c
// BarrierMegacore — 0x1d5222a0 (gate, decompiled)
v8 = target;                                         // region()->...->target
if (TpuChipConfig::Megacore(chip_config(target))) {  // Megacore() mode
    if (*(int*)(chip_config + 124) >= 2) {           // TpuChipConfig+0x7c >= 2  (cores per megachip)
        v9 = Target::CoresPerChip(target, 0);        // TC cores
        if ((int)v9 >= 2) {
            CHECK(!region()->InPrimaryOrSecondaryRegion());  // llo_region_builder.cc:1751
            ...                                       // fold body
        }
    }
}
// otherwise: no barrier emitted

TpuChipConfig::Megacore() — 巨核模式判别式(偏移路径 target[+0x3b8] → +0x18 → Megacore())。
TpuChipConfig+0x7c ≥ 2 — 每兆芯片核心数字段(已确认偏移；≥ 2 测试是兆核条件)。
CoresPerChip(kTensorCore) == 2 — 在 llo_region_builder.cc:1772(标量路径)和 :1788(矢量路径)处重新检查为 CHECK(core_count == 2)；如果一个超级核心芯片报告除了两个 TensorCore 以外的任何信息，该函数就会中止。这是偶/奇对不变量。
CHECK(!region()->InPrimaryOrSecondaryRegion())(llo_region_builder.cc:1751，VLOG "BarrierMegacore in primary or secondary region, hlo: ") - 势垒不得在两个每核子区域之一内部发射；它在连接它们的折叠点处发射。

2.2 指纹屏障id

与 BarrierConfig(带有彩色 id)不同，超级核心屏障 id 是在发出时构建的每指令指纹 - 在两个核心上稳定，因此双方计算相同的 SFLAG cookie：

// BarrierMegacore — 0x1d5222a0 (id construction, decompiled)
v19 = Fingerprint2011(operand_ptr, ...);                 // hash of the region's operand
if (hlo) {                                                // HloInstruction*
    FormatPack(&buf, "%d", /*one arg=*/2);                // "2" -> a literal discriminator
    v24 = Fingerprint2011(buf, ...);
    v19 = FingerprintCat2011(v19, v24);
    HloInstruction::unique_id(hlo);                       // folded into the str-format args
}
// per-module monotone cookie:
v27 = *(int*)(module + 68);                               // cookie_ordinal_
CHECK(v27 < INT32_MAX, "cookie_ordinal_ < std::numeric_limits<int32_t>::max()");  // llo_module.h:147
*(int*)(module + 68) = v27 + 1;                           // post-increment
v28 = FingerprintCat2011(v19, v27);
cookie = (uint32_t)(v28 ^ (v28 >> 32));                   // fold to 32 bits
// annotation: "BarrierMegacore cookie: %u @ %s:%d"
```text

该 cookie 注释为 `"BarrierMegacore cookie: %u @ %s:%d"`(使用 `absl::SourceLocation` 文件/行)，因此两个核心的屏障在转储中可被识别为匹配对。因此，屏障 ID 来自 `HloInstruction::unique_id()` + 每个模块 `cookie_ordinal_`，**不是**来自 `BarrierConfig.id()`。

### 2.3 偶/奇同步相加折叠

折叠本身：加载超级核心 SFLAG 插槽，然后每个合作伙伴发出一个远程同步标志添加到*其他*核心，并等待合作伙伴的碰撞。

```c
// BarrierMegacore — 0x1d5222a0 (fold, core_count == 2, decompiled & condensed)
chip   = LloRegionBuilder::ChipId(this, 0);
core   = LloRegionBuilder::CoreIndex(this);               // this core's index (even/odd)
n      = Target::GetMegacoreBarrierSyncFlagNumber(target); // = base+count  (reserved megacore slot)
sflag  = LloRegionBuilder::SflagImmPtr(this, n, "megacore barrier sync flag", 26);

for (v62 = 1; v62 != core_count /*==2*/; ++v62) {         // one partner when core_count==2
    delta   = ScalarU32Constant(v62);                     // +1
    partner = SaddS32(core, delta);                       // core + 1
    partner = SmodU32(partner, core_count /*==2*/);        // (core + 1) mod 2  ->  the OTHER core
    addr    = EncodeRemoteSyncFlagAddress(this, sflag, &chip, /*multicast=*/0);
    op      = LloInstruction::CreateVectorSyncFlagAddRemote(addr, ScalarU32Constant(1), region, ...);
    LloRegion::AppendInstruction(region, op, 0);          // bump partner's megacore SFLAG by 1
}
wait = VwaitGeSV(this, sflag, ScalarU32Constant(core_count - 1) /*==1*/, 0);  // wait for partner's bump
wait.set_annotation_if_not_constant("barrier-megacore-wait");
VsyncAdd(this, sflag, ScalarU32Constant(1 - core_count) /*== -1*/);          // decrement back to 0
// ... ScheckEq verification on the slot ...

算术 SmodU32(CoreIndex + step, core_count) 是偶数/奇数合作伙伴映射：对于 core_count == 2，它发送到 (core + 1) mod 2，即核心 0 向核心 1 发出信号，核心 1 向核心 0 发出信号。每个核心将合作伙伴的保留超级核心 SFLAG 副本增加 1 (VectorSyncFlagAddRemote，通过 EncodeRemoteSyncFlagAddress 寻址)，等待直到其自己的槽到达 core_count − 1 == 1(VwaitGeSV，注释为 "barrier-megacore-wait")，然后递减 1 − core_count == −1(VsyncAdd)以使槽保留为零以供重用。 MegacoreBarrierAnn 注释变体附加到控制流推送 (VCcfPush / SCcfPush)，以便稍后的传递可以识别折叠。

注意 — 页面发出两个几乎相同的主体：一个向量路径(VectorU32Constant → VCcfPush → … → VCcfPop → SimplifyVtos，在主 core_count == 2 臂上获取)和一个标量路径(ScalarU32ConstantImpl → SCcfPush，在预测字节选择它时获取)。两个门都在 core_count == 2 上，并且都运行相同的 GetMegacoreBarrierSyncFlagNumber → SflagImmPtr("megacore barrier sync flag") → SaddS32/SmodU32 → VsyncAddRemote 折叠；区别在于矢量与标量控制流推送。详细的 VsyncAdd / VwaitGeSV SFLAG 握手语义位于屏障中。

2.4 程序描述符状态折叠(运行时孪生)

net_util::SynchronizeProgramDescriptorStatesMegacore (0x1c697540) 是兆核 SFLAG 插槽的第二个消费者。它会合两个内核的程序描述符状态字(SMEM 控制字)而不是数据屏障，并且门控相同：

// SynchronizeProgramDescriptorStatesMegacore — 0x1c697540 (decompiled, condensed)
if (!TpuChipConfig::Megacore(chip_config) || *(int*)(chip_config + 124) < 2)
    return;                                              // same Megacore() + 0x7c>=2 gate
// fingerprint id: unique_id + "%d"(2) + ToShortString + cookie_ordinal_ (++)
//   annotation "SynchronizeProgramDescriptorStatesMegacore cookie: %u @ %s:%d"
CHECK(CoresPerChip(0) / LogicalDevicesPerChip(0) == 2,
      "sync_b.target().TensorCoresPerLogicalDevice() == 2");   // net_util.cc:1337
n     = Target::GetMegacoreBarrierSyncFlagNumber(target);      // base+count
sflag = SflagImmPtr(b, n, "megacore barrier sync flag", 26);
pds   = SmemWordImmPtr(b, Target::ProgramDescriptorStateWordOffset(target),
                       "program descriptor state", 24);
// PrimaryCoreRegionBuilder + SecondaryCore -> EnqueueRemoteSst(pds) to the partner;
VwaitGeSV(b, sflag, 1, 0).set_annotation_if_not_constant("...wait");
VsyncAdd(b, sflag, SimmS32(-1));
Sfence(b);
// ScheckEq verification
```text

CHECK `TensorCoresPerLogicalDevice() == 2`(`= CoresPerChip(TC) / LogicalDevicesPerChip(TC)`、`net_util.cc:1337`)是 **每个逻辑设备两个核心** 巨核不变量的第三个独立语句。此折叠使用 `EnqueueRemoteSst` 将主核心的程序描述符状态复制到辅助核心，并由相同的保留巨核 SFLAG 插槽隔离。

### 2.5 megacore SFLAG 插槽恰好有 3 个消费者

`Target::GetMegacoreBarrierSyncFlagNumber` (`0x1d60f4e0`) 的完整 `.text` `E8/E9` 交叉引用恰好找到三个调用者 - 确认保留的 `base+count` 槽仅由运行时/分析路径消耗，而不是通过 `BarrierConfig`：

| 消费者 (VMA) | 角色 |
|---|---|
| `LloRegionBuilder::BarrierMegacore` (`0x1d522641`) | TC 偶数/奇数数据屏障(§2.3) |
| `net_util::SynchronizeProgramDescriptorStatesMegacore` (`0x1c6977fc`) | 程序描述符状态折叠(§2.4) |
| `RaceAnalyzerStepper::PreProcessEvent` (`0x10bb3229`) | 在 LLO 竞赛分析器中保留/记入槽位 |

---

## 3. `GetMegacoreBarrierSyncFlagNumber` — 预留插槽访问器

超级核心 SFLAG 编号由 `Target::GetMegacoreBarrierSyncFlagNumber` (`0x1d60f4e0`) 计算为 `base + count` — 即可用的 per-id 屏障窗口上方的**第一个插槽：

```c
// Target::GetMegacoreBarrierSyncFlagNumber — 0x1d60f4e0 (decompiled)
__int64 Target::GetMegacoreBarrierSyncFlagNumber(Target *this) {
  if (!TpuChipConfig::Megacore(chip_config(this)))
    CHECK_FAIL("topology_->chip_config().Megacore()");          // target.cc:154
  return (uint32_t)(this[560] + this[561]);                     // base(+0x8c0) + count(+0x8c4)
}

this[560] 和 this[561] 是 _DWORD 索引，即字节偏移量 0x8c0(基数)和 0x8c4(计数)。访问器是 Megacore() 门控，具有硬 CHECK (target.cc:154) — 在非兆核芯片上调用它会中止。对比索引同一块顶部的两个同级保留访问器：

访问器 (VMA)	公式	插槽	消费者
`GetMegacoreBarrierSyncFlagNumber` (`0x1d60f4e0`)	`base + count`	`base+count+0`	`BarrierMegacore` / `SyncProgDesc…Megacore` / `RaceAnalyzer` (§2.5)
`GetAllReduceSyncFlagNumber(n)` (`0x1d60f440`)	`base + count + n + 1` (`0 < n < 3`)	`+2`, `+3`	`RotatedPincer` / `AsyncPincer` `InitSyncFlags` (分层/钳形)
`GetGlobalBarrierSyncFlagNumber` (`0x1d60f420`)	`base + count + 4`	`+4`	`GetBarrierSyncFlag(GLOBAL)` + 树障

GetAllReduceSyncFlagNumber 检查其相位参数 phase > 0 (target.cc:143) 和 phase < 3 (target.cc:144)，因此 n ∈ {1, 2} 映射到 base+count+2 和base+count+3; base+count+1 插槽(非法的 n=0)是一个未使用的间隙。这四个占用的槽位加上间隙正好是Target::Init(§4)保留的五个槽位。

4. TC SFLAG 块初始化(`Target::Init`)

Target::Init(0x1d60fc20、target.cc)在启动时从芯片的专用 SFLAG 范围填充 TensorCore 保留屏障块。这两个字段存储在反编译中进行了字节确认：

// Target::Init — 0x1d60fc20 (decompiled, the TC barrier-block filler)
SpecialPurposeSyncFlags = TpuChipConfig::GetSpecialPurposeSyncFlags(chip_config /* kTensorCore */);
if (!SpecialPurposeSyncFlags)
    DieBecauseNull("chip_config.GetSpecialPurposeSyncFlags(::tpu::TpuCoreType::kTensorCore)");  // target.cc
// ... copy + contiguity-check the int range (must be a contiguous ascending run) ...
*((_DWORD *)target + 560) = *v286;        // Target+0x8c0 = arr[0]           (TC base)
*((_DWORD *)target + 561) = v284 - 5;     // Target+0x8c4 = arr.size() - 5   (TC usable count)
```text

- **`GetSpecialPurposeSyncFlags(kTensorCore)`** (`0x20afcf40`) 返回 SFLAG 数字的 `{ptr, count}` 向量 — 芯片配置表 (`TpuChipConfig+0x2a0 + core`) 的每个 `TpuCoreType` 切片。 `Target::Init` 复制它，然后断言(展开的 `inc/cmp/jne` 循环)这些数字形成**连续的升序整数范围**，以便 `base + offset` 寻址有效。
- **`Target+0x8c0 = arr[0]`** — 第一个特殊用途的 SFLAG 编号(块基数)。
- **`Target+0x8c4 = arr.size() − 5`** — 每个 id 可用的屏障计数。 **`−5` 为指定的运行时访问器保留该范围的前五个插槽**：

```text
TpuChipConfig::GetSpecialPurposeSyncFlags(kTensorCore)  =  [ base ............................. base+size-1 ]
                                                              |<--- usable per-id window --->|<-- 5 reserved -->|
   Target+0x8c0 = base
   Target+0x8c4 = size - 5  =  count            base .. base+count-1   :  REPLICA(2)/CUSTOM(3) ids (id < count)
                                                base+count+0           :  MEGACORE   (GetMegacoreBarrierSyncFlagNumber)
                                                base+count+1           :  (unused gap; GetAllReduceSyncFlagNumber(0) illegal)
                                                base+count+2           :  AllReduce phase 1
                                                base+count+3           :  AllReduce phase 2
                                                base+count+4           :  GLOBAL     (GetGlobalBarrierSyncFlagNumber)

所以 megacore 屏障槽是第一个保留槽 (base+count+0)，位于 REPLICA/CUSTOM 从中绘制的可用 per-id 窗口的正上方。在 Init 安装实际的每代值之前，基本 Target 构造函数 (0x1d615340) 将块(+0x8c0 = 0、+0x8c4 = 0)设置为零。

注意 — 每个芯片代号的文字 GetSpecialPurposeSyncFlags(TC).size()(因此具体的 Target+0x8c4 计数)是从 chip_config / chip_parts 原型填充的，而不是从二进制文件中提取的； writer 及其公式(base = arr[0]、count = size − 5)是字节确认的(仅针对每个代号文字为 LOW)。 SparseCore 镜像(SparseCoreTarget::Init、基础 +0x1d0、计数 +0x1d4 和无 −5)和三区域 SFLAG 分区位于 SpecialPurposeSyncFlags 中 — SC 引擎在顶部不保留任何内容，因为它的屏障在块内 * 带有保留 id。

4.1 为什么巨核插槽是`base+count`，而不是`base+id`

这就是MEGACORE(4)故意绕过BarrierConfig的症结所在。每个集合的 REPLICA/CUSTOM 屏障位于 [base, base+count) 窗口内部，由彩色 id (§1.1) 索引。超级核心折叠不是每个集体 - 它是两个核心之间的固定芯片级交汇处，因此它需要一个单个专用插槽，每个ID的屏障都不会发生冲突。将其放置在 base+count(可用窗口上方)并通过 GetMegacoreBarrierSyncFlagNumber(而不是通过 GetBarrierSyncFlag)到达它可以保证这一点。如果 MEGACORE(4) 是真正的 BarrierConfig，它会解析为 base + id(在每个集合窗口内)，这正是设计要避免的冲突。

5. 与更高级别 `MegacoreFusion` HLO 通证的关系

上面的运行时折叠与编译器级 xla::MegacoreFusion HLO 传递(RunImpl 0x110d8f00、DoMegacoreFusion 0x110d8860、TU tpu_megacore_fusion.cc)不同。该过程在“降低”之前运行：它将全归约与计算操作配对(例如 FindARPair 0x110d5560、FindConvARsMatch 0x110d2980、GetAllReduceCosts 0x110d1bc0)，以便一个核心的集合与另一个核心的 matmul/卷积重叠 - 调度跨超级核心对的融合。本页上的 BarrierMegacore / SynchronizeProgramDescriptorStatesMegacore 构造是“降低时间”交会点，使此类配对时间表正确：它们迫使两个核心在折叠点重新达成一致。重新实现者必须保持两层分离——HLO 传递决定“什么”重叠，运行时折叠保证两半在保留的超级核心 SFLAG 上“同步”。

注意 — MegacoreFusion HLO pass 内部结构(AR/conv 配对启发式、延迟限制匹配、它查询的成本模型)超出了此处的范围；此页面仅记录它是一个单独的较早层。它可以配对的钳形/二项式 AllReduce 发射器是分层/钳形和二项式/递归加倍。

6. 重新实施清单

不发出 MEGACORE(4)。 无传递写入 BarrierConfig.barrier_type == 4。将配置生成器限制为 GLOBAL(1) / REPLICA(2) / CUSTOM(3)；使 0/4 无法访问。在 GetBarrierSyncFlag 中，类型 {2,3,4} 共享一个 base + id 臂绑定检查 id < count — 没有每种类型的 2/3/4 逻辑。
以三种方式控制运行时折叠。 仅当 TpuChipConfig::Megacore() ∧ TpuChipConfig+0x7c ≥ 2 ∧ CoresPerChip(TC) == 2 时才发出 BarrierMegacore。重新断言 CHECK(core_count == 2)(以及描述符折叠的 TensorCoresPerLogicalDevice() == 2)—偶数/奇数折叠仅针对每个逻辑设备的两个内核定义。
从指纹而不是配置 ID 构建屏障 id。 Fingerprint2011(operand) ⊕ 鉴别器的 "%d" 格式 + HloInstruction::unique_id() + 每个模块的单调 cookie_ordinal_ (检查 < INT32_MAX，后增量)，折叠为32位。注释 "BarrierMegacore cookie: %u @ %s:%d"。
运行偶/奇同步添加。 每个合作伙伴步骤 v ∈ [1, core_count)：partner = SmodU32(SaddS32(CoreIndex, v), core_count)； VectorSyncFlagAddRemote(EncodeRemoteSyncFlagAddress(megacore_sflag, chip), +1)。然后将VwaitGeSV(megacore_sflag, core_count − 1)(注"barrier-megacore-wait")和VsyncAdd(megacore_sflag, 1 − core_count)重置为零。
预留兆核插槽。 在 Target::Init、base = arr[0] → +0x8c0、count = GetSpecialPurposeSyncFlags(TC).size() − 5 → +0x8c4 中；连续性检查范围的前五个是 {Megacore @base+count, gap @+1, AllReduce1 @+2, AllReduce2 @+3, Global @+4}。通过 Megacore() 门控访问器将兆核插槽计算为 base + count — 绝不是 base + id。
也同步程序描述符状态。 SynchronizeProgramDescriptorStatesMegacore 将主核心的 ProgramDescriptorStateWordOffset SMEM 字复制到辅助核心 (EnqueueRemoteSst)，并在同一保留的超级核心 SFLAG 上进行防护 (Sfence)。

交叉引用

集体概述 — 两个基板集体堆栈和策略选择器； megacore 是与每轴环计划正交的每芯片折叠。
AllReduce 分层/钳形 — 双向环融合，其臂 MegacoreFusion HLO 通道可以跨两个核心配对；也是保留的 AllReduce SFLAG 插槽 (base+count+2/+3) 的使用者。
二项式/递归加倍 — 自完成蝶形发射器(延迟限制)，另一个 AllReduce 系列。
物理核心放置 — 逻辑设备 ID 如何映射到物理 TensorCore(此折叠同步的偶数/奇数对)。
ReduceScatter — 在超级核心折叠之前针对每个核心运行的减少分散阶段。
Megacore 偶数/奇数分割 — 扭曲拓扑：哪个核心是主核心，哪个核心是辅助核心，以及该对如何映射到环面。
屏障 — 基于 SFLAG 的屏障模型和完整的 BarrierType 分类/握手协议。
TensorCore 屏障 — TensorCore 侧势垒降低和 SFLAG 绑定。
SpecialPurposeSyncFlags — Target::Init / SparseCoreTarget::Init SFLAG 块填充程序和三区域分区。
InferBarrierConfig — BarrierConfig 标准化器(两种类型生产者之一，确认永远不会写入 4)。
返回索引

巨核集体融合 ​

摘要 ​

1. MEGACORE(4) 标记绝不是 BarrierConfig ​

1.1 GetBarrierSyncFlag的死4型臂 ​

2.2 指纹屏障id ​

2.4 程序描述符状态折叠(运行时孪生) ​

4. TC SFLAG 块初始化(Target::Init) ​

4.1 为什么巨核插槽是base+count，而不是base+id ​

5. 与更高级别 MegacoreFusion HLO 通证的关系 ​

6. 重新实施清单 ​

交叉引用 ​