屏障着色

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

摘要

BarrierColoring is the TensorCore (TC) barrier-deduplication engine: a greedy graph-coloring pass that decides which concurrent TC collectives may share a barrier sync flag and which must be forced onto distinct ones. SparseCore 屏障传递纯粹在静态环配置哈希上进行重复数据删除(没有调度重叠的概念)，而 TC 传递在异步集合的“生命范围”上构建一个“干扰图”并对其进行着色：两个共享 TensorCoreBarrierKey 但其异步 start..done 窗口在调用图排序调度中重叠的集合会获得干扰边缘并被着色。概述描述了基于 SFLAG 的屏障模型和 BarrierType 枚举； SFLAG号码绑定描述了编译器屏障 → 硬件 SFLAG 编号映射； InferBarrierConfig 描述了每代 SFLAG 映射源。 This page owns the greedy coloring engine (the Run → AssignColorForConflictingOp → VisitNodes/CollectConflictInfo → AssignColorsGreedy pipeline, the conflict-edge construction, and the first-fit color scan) plus the BarrierConfig → chip-SFLAG lowering (CustomKernelEmitter::Emit → MaybeInsertGlobalBarrier / RunPasses) 将彩色屏障转换为具体的同步标志 memref。

The engine is templated as xla::jellyfish::BarrierColoring<Policy<TensorCoreBarrierKey>> and instantiated twice — once for collective-permute (AsyncCollectivePermutePolicy) and once for async-barrier'd all-to-all (AsyncAllToAllWithAsyncBarrierPolicy). The two instantiations share all algorithm code and differ only in six policy predicates (what counts as an async start/done, how to derive the key).每个通道的输出是一对：一个 map<HloInstruction*, long> 为每个 TC 集合分配一个颜色，以及一个采用非零颜色的操作的 flat_hash_set<HloInstruction*> - “必须获得一个不同的屏障”集，它为障碍分配生产者中的 DetermineBarrierConfigForKey 的 has_conflict 参数提供数据(@ 0x109c6fa0)。

重新实现，合约为：

Dedup 是图形着色，而不是散列。 颜色 0 是默认/共享屏障；颜色 >= 1 是冲突造成的额外明显障碍。密钥的干扰图的色数等于该密钥的最大并发运行中集合，这正是该密钥消耗的不同屏障 id 的数量。
冲突谓词是生存范围重叠。 两个异步集合发生冲突，当且仅当它们的 async-start … async-done 窗口按计划顺序重叠时。调度必须通过 CallGraph DFS(不是无序迭代)遍历，以便构建器可以跟踪每个程序点当前打开的集合。
颜色搜索是确定性的首次拟合。 对于每个节点，构建邻居颜色集，然后扫描 0, 1, 2, … 并获取邻居未使用的第一个整数(最小可用颜色)。确定性输入顺序(std::map)+确定性搜索⇒可重现的着色，这对于缓存稳定编译很重要。
**决策和降低仅通过原型连接。**所选的 BarrierConfig {type, id} 通过分配传递写入 HLO 的 BackendConfig 并在发出时由 CustomKernelEmitter::Emit 读回。着色和降低之间没有共享的内存状态。
两个屏障臂都以相同的 SFLAG 原语结束。 type 1 全局 (GLOBAL) 屏障降低到 AllocateAtOffsetOp(MemorySpace::sflag)，包裹在 tpu.sem_signal/tpu.sem_wait(TC 树)上的 scf.for 循环中屏障)； type 2/3 每键 (REPLICA/CUSTOM) 屏障流经 RunPasses 到 SparseCore 每环屏障使用的相同 GetSyncFlagForBarrierId → AllocateAtOffsetOp(sflag) 路径。同一芯片SFLAG块；不同的原子操作家族。 type 是降低比较的文字 BarrierConfig.type 原始整数(== 3 ⇒ 每个键)；枚举名称为 GLOBAL=1/REPLICA=2/CUSTOM=3(请参阅概述)。


发动机类别	`xla::jellyfish::BarrierColoring<Policy<TensorCoreBarrierKey>>`(TU `barrier_assignment.cc`)
ACP `Run`	`@ 0x109cf600`(集体置换策略)
A2A `Run`	`@ 0x109d1a60`(异步屏障全面策略)
编排器	`AssignColorForConflictingOp` `@ 0x109cf6c0` (ACP) / `@ 0x109d1b20` (A2A)
日程步行	`VisitNodes` `@ 0x109d1240` → `VisitNodesInternal` `@ 0x109d1420`(CallGraph DFS)
干扰生成器	`CollectConflictInfo` `@ 0x109d4a20` (ACP) / `@ 0x109d5a60` (A2A)
颜色搜索	`AssignColorsGreedy` `@ 0x109d0c80` + 内部 lambda `@ 0x109d0de0`(首次适配)
冲突节点	`ConflictInfo`，`sizeof == 0xb0`，分配`@ 0x109cfd56`
SFLAG 降低	`CustomKernelEmitter::Emit` `@ 0x1321ad60` → `MaybeInsertGlobalBarrier` `@ 0x1321ac20` / `RunPasses` `@ 0x13202780`
全局 SFLAG 原语	`sparse_core::AllocateAtOffsetOp::create` `@ 0x145a5aa0` (`MemorySpace::sflag`)
TC 屏障操作	`tpu::SemaphoreSignalOp` `@ 0x14b442e0` / `tpu::SemaphoreWaitOp` `@ 0x14b45460`
信心	已确认(引擎 + 降低机构已反编译和双重检查)，除非一行另有说明

引擎类及其两个实例化

用途

BarrierColoring 的存在是因为共享相同 TensorCoreBarrierKey 的两个 TC 集合“不”自动安全地共享屏障同步标志。如果两者同时运行(例如，双缓冲集体置换管道，其中一个 CP 的 start..done 窗口打开，而同一密钥的前一个 CP 仍处于活动状态)，则它们必须在不同的同步标志上发出信号，否则会破坏彼此的完成计数。引擎通过将每个异步集合建模为一个间隔(start 打开，done 关闭)并对生成的干扰图进行着色来准确检测这种情况。

该类是一个函数对象模板：

cpp

// xla::jellyfish (TU platforms/xla/service/jellyfish/barrier_assignment.cc)
template <typename Policy>   // Policy = AsyncXxxPolicy<TensorCoreBarrierKey>
class BarrierColoring {
 public:
  // returns {color map, conflict-op set}
  StatusOr<std::pair<std::map<HloInstruction*, long>,
                     flat_hash_set<HloInstruction*>>>
  Run(HloModule* module);

 private:
  StatusOr<...> AssignColorForConflictingOp(
      HloModule* module, const flat_hash_set<string_view>& computation_filter);
  void VisitNodes(HloModule*, FunctionRef<Status(HloInstruction*, ConflictMap&)>,
                  ConflictMap*);
  Status CollectConflictInfo(HloModule*, ConflictMap&, HloInstruction*);
  std::map<HloInstruction*, long> AssignColorsGreedy(
      const std::map<HloInstruction*, std::set<HloInstruction*>>& adjacency);
};
```text

### 两项政策

这两个实例化是通过 6 槽 vtable 实现的。仅谓词不同；该算法是字节相同的(每个函数的 A2A 变体 - 例如 `AssignColorsGreedy @ 0x109d30e0`、`CollectConflictInfo @ 0x109d5a60` - 与 ACP 变体具有相同的堆栈帧和相同的循环)。

| vtable插槽 | 方法 | `AsyncCollectivePermutePolicy` | `AsyncAllToAllWithAsyncBarrierPolicy` |
|---|---|---|---|
| `+0x10` | `IsAsyncStart` | `opcode == 0x24`(CP-启动) `@ 0x109cc080` | `IsCustomCall(kBarrierStart)` `@ 0x109cc640` |
| `+0x18` | `IsAsyncDone` | `opcode == 0x23`(CP-完成)`@ 0x109cc0a0` | `opcode == 0x11 && GetBarrierStartFromAsyncStart != 0` `@ 0x109cc660` |
| `+0x20` | `GetKey`(构建密钥) | 正常化→启动`@ 0x109d3ec0` | `async_wrapped_instruction` → 钥匙 `@ 0x109d4460` |
| `+0x28` | `GetStartForDone` | `done->operand(0)` `@ 0x109cc0e0` | `GetBarrierStartFromAsyncStart(done)` `@ 0x109cc680` |
| `+0x30` | `GetCollectiveForStart` | 身份`@ 0x109cc0c0` | `@ 0x109d4400` |
| `+0x38` | `ShouldForceGlobalBarrier` | 简并-CP门`@ 0x109d3f20` | 底座 `false` `@ 0x109cac00`(继承) |

> **注意 - 已验证的操作码。** ACP 的反编译 `IsAsyncStart` 主体实际上是 `return *(_BYTE *)(op + 0xc) == 36;`(36 = `0x24`，集体排列启动)，而 `IsAsyncDone` 是 `== 35;` (`0x23`，集体排列完成)。异步屏障助手是 `async_barrier_util::kBarrierStart @ 0xabe9620`、`GetBarrierStartFromAsyncStart @ 0x11007c80`(HLO 异步启动 `0x11`、后端类型 `== 2`、`operand(0)` 操作码 `0x81` = 屏障启动)和 `GetAsyncStartFromBarrierStart @ 0x11007d20`。

这两个策略需要*单独的*传递，因为它们的异步开始/完成定义是不相交的：HLO Collective-permute-start/done 操作码上的 CP 路径键直接，而 `all-to-all` 上的 A2A 路径键封装在 HLO 异步开始/完成内的异步屏障自定义调用 (`kBarrierStart`) 中对。异步屏障完整操作系列记录在 [张量核势垒](tensorcore-barrier.md) 页面上。

---

## `Run` — 薄包装

`Run` (`@ 0x109cf600` ACP / `@ 0x109d1a60` A2A) 是样板：它对空 `flat_hash_set<string_view>` 进行零初始化(“限制这些计算的着色”名称过滤器 - 此处的空意味着*所有* TC 线程计算)，通过`raw_hash_set` 构造函数 `@ 0x10912d00`，然后尾部调用 `AssignColorForConflictingOp(module, &name_set) @ 0x109cf6c0`。返回时，它释放集合的支持数组 (`DeallocateBackingArray @ 0x21118960`) 并返回协调器在 `*this` 中构建的 `StatusOr<pair<map<HloInstruction*,long>, flat_hash_set<HloInstruction*>>>`。该算法完全在`AssignColorForConflictingOp`中。

---

## `AssignColorForConflictingOp` — 协调器

`AssignColorForConflictingOp`(`@ 0x109cf6c0` ACP / `@ 0x109d1b20` A2A)接收`(this = result, module, name_set)`并运行五个阶段。

### 1 — 枚举计算

`HloModule::computations(name_set) @ 0x10944b40` 生成 TC 线程计算视图，复制到堆栈帧上。

### 2 — 按键分组到 `ConflictInfo` (`@ 0x109cf810 … 0x109d0404`)

遍历每个计算的每条指令。对于每个操作，构建其 `TensorCoreBarrierKey` 内联(读取操作码；通过 `ReplicaGroup` ctor `@ 0x20e43de0` 从 `[op+0xd0]`/`[op+0xd8]` 复制副本组；复制 src-tgt 对；通道奇偶校验)，然后将 `__emplace` 转换为 `std::map<TensorCoreBarrierKey, ConflictInfo>`植根于框架 `-0x108`。在 **新** 密钥上，堆分配 `ConflictInfo` 节点 (`mov edi, 0xb0; call _Znwm` `@ 0x109cfd56`) 并将密钥的副本组 / src-tgt 向量复制到其中 - 该节点携带自己的密钥有效负载副本，因此它在密钥的生命周期内存活。该地图由 `TensorCoreBarrierKey::operator< @ 0x109d6620` 订购。每个 `ConflictInfo` 从而对共享密钥的所有集体进行分组。

`0xb0` 字节 `ConflictInfo` 节点子布局(由构造存储 `@ 0x109cfd5b … 0x109cff7b` 证明的偏移；从使用中推断出的子成员*语义*)：

| 偏移 | 会员 |
|---|---|
| `+0x10` | 地图节点键后向指针 |
| `+0x20` | 密钥操作码字节镜像 |
| `+0x28 … +0x38` | 副本组向量(已复制) |
| `+0x40 … +0x50` | src-tgt-对向量(已复制) |
| `+0x58 … +0x78` | 复制钥匙尾 |
| `+0x80` | 每键邻接 `map<HloInstruction*, set<HloInstruction*>>` 根 |
| `+0x98` | 每键彩色 `map<HloInstruction*, long>` 根 |

### 3 — 构建干涉图 (`@ 0x109d0469`)

设置`CollectConflictInfo`回调(`FunctionRef`；`InvokeObject @ 0x109d4a00`)并调用`VisitNodes(module, callback, &conflict_map) @ 0x109d1240`。这会使用每个键的干扰边填充节点 `+0x80` 处的每个 `ConflictInfo` 的邻接图。

### 4 — 每键贪婪着色 (`@ 0x109d04dd … 0x109d066c`)

对于每个 `ConflictInfo` 条目(映射遍历，`node+0x80` 处的邻接映射根)，调用 `AssignColorsGreedy(this, &out_colors, &node->adjacency) @ 0x109d0c80`，然后将返回的每键着色通过 `__insert_range_unique @ 0x109cc380` 合并到模块范围的颜色映射到框架 `-0x2b8` 树中。

### 5 — 发出冲突操作集 (`@ 0x109d0885 … 0x109d0880`)

对于分配颜色为 `> 0` 的操作(它们不能采用颜色 0 = 共享/默认屏障，即它们与相同密钥对等点冲突)，将操作插入到返回的 `flat_hash_set<HloInstruction*>` (`find_or_prepare_insert_large @ 0xe63dbe0`) 中。

> **NOTE — `RET_CHECK`.** A `RET_CHECK` at line `0xdf` (`@ 0x109d05f6`) guards the orchestrator against a malformed coloring intermediate.

输出为 `(color map, conflict-op set)`。 [障碍分配生产者](overview.md) 消耗两者：冲突集提供 `DetermineBarrierConfigForKey @ 0x109c6fa0` 的 `has_conflict` 参数，这正是图冲突如何将共享的 `REPLICA(2)` 屏障转变为新的 `CUSTOM(3)` 屏障的方式 - 请参阅 [着色方式 `BarrierConfig`](#how-the-coloring-feeds-barrierconfig)。

---

## `VisitNodes` / `VisitNodesInternal` — 日程步行

`VisitNodes @ 0x109d1240`：

1. `CallGraph::Build(module, name_set) @ 0x1e5579e0` — 通过 TC 线程计算构建调用图。
2. `CallGraph::GetNode(entry_computation) @ 0x1e556600` — 入口节点。
3. `VisitNodesInternal(module, callback, call_graph, entry_node, visited_map, conflict_map) @ 0x109d1420`(A2A 的 `@ 0x109d3880`) - 由访问过的戳记的 `flat_hash_map<CallGraphNode*, long>` 键入的递归 CallGraph DFS，**按计划顺序**遍历每个计算的指令并调用每个指令的回调。

The traversal order is essential: `CollectConflictInfo` detects conflicts by tracking which async collectives are *currently open* at each program point, so it must see instructions in schedule order — hence the CallGraph DFS rather than an unordered map iteration.

---

## `CollectConflictInfo` — 冲突谓词(干扰图生成器)

`CollectConflictInfo` (`@ 0x109d4a20` ACP / `@ 0x109d5a60` A2A) 具有签名 `lambda(HloInstruction* op, map<TensorCoreBarrierKey, ConflictInfo>& conflict_map) → Status`。它使用上表中的插槽(`IsAsyncStart [vt+0x10]` `@ 0x109d4a4c`、`IsAsyncDone [vt+0x18]` `@ 0x109d4a62`、`GetStartForDone [vt+0x28]` `@ 0x109d4ac4`、`GetCollectiveForStart [vt+0x30]`)通过虚表 `[op_policy]` 达到策略`@ 0x109d4b01`)。

每个访问过的操作：

- **`IsAsyncStart(op)`** → 操作 **打开** 异步集体活动范围。计算其密钥(策略`GetKey`链)，在`conflict_map`(`TensorCoreBarrierKey::operator< @ 0x109d6620`树遍历)中查找其`ConflictInfo`。将 `op` 添加到 **实时异步集** — 本地 SwissTable，通过基于 crc32 的 `MixingHashState` 路径(`_mm_crc32_u64`、`@ 0x109d4c7b … 0x109d4cf3`)在 `HloInstruction*` 指针上散列成员资格。
- For every **other** currently-live async collective in the set: add a **bidirectional interference edge** — insert each into the other's `ConflictInfo` adjacency `set<HloInstruction*>` (`HloPtrComparator`-ordered set insert, `@ 0x1e5a7b20`, via `__emplace_unique`). `start..done` 范围重叠的两个异步集合因此**冲突**并且不能共享屏障同步标志。
- **`IsAsyncDone(op)`** → `GetStartForDone(op) [vt+0x28]` → **关闭**该开始的现场范围(从现场设置中删除)。
- **`ShouldForceGlobalBarrier(op) [vt+0x38]`** — 当为 true 时(ACP 退化集体置换情况 `@ 0x109d3f20`)，操作将转向全局屏障，无论颜色如何。
- A `RET_CHECK`(`MakeErrorStream @ 0x20cf6b80`、`add_ret_check_failure @ 0x20cf6be0`，行 `0x9a` `@ 0x109d4b84`)在格式错误的异步对上触发(没有匹配开始的完成，反之亦然)。

```text
schedule order →  ... CP-start(A) ... CP-start(B) ... CP-done(A) ... CP-done(B) ...
live set:         { }   {A}            {A,B}          {B}            { }
edges added:                 A—B   (A and B overlap → interference edge)

这是与 SparseCore 屏障传递的结构差异：SC 通过 FlatHashMap 在静态环配置密钥上进行重复数据删除，没有计划重叠的概念； TC 在实时异步集体干扰图上传递重复数据，因此即使具有相同 TensorCoreBarrierKey 的两个集体如果同时飞行，也会被分开。

`AssignColorsGreedy` — 首次适配最小可用颜色

AssignColorsGreedy (@ 0x109d0c80 ACP / @ 0x109d30e0 A2A) 采用一个 ConflictInfo 的干涉图 — const std::map<HloInstruction*, set<HloInstruction*>>& adjacency — 并返回 std::map<HloInstruction*, long>(节点 → 颜色)， HloPtrComparator-订购。

外循环 (@ 0x109d0c80)：迭代邻接图； for each node take its neighbor set (rbx+0x28) and call the inner lambda @ 0x109d0de0, then insert (node_ptr = [rbx+0x20], color) into the output color map via __try_key_extraction_impl @ 0x109d47e0 / __emplace_unique.

内部 lambda (@ 0x109d0de0) — 最小自由颜色搜索，逐字节确认：

构建本地flat_hash_set<long> used(raw_hash_set，策略@ 0x21616168)。
For each neighbor in the neighbor set (walked in HloPtrComparator order): look the neighbor up in the so-far output color map (__try_key_extraction_impl); if it already has a color, insert that color into used — the SwissTable insert uses _mm_crc32_u64 to hash the long color (@ 0x109d0e6e region), PrepareInsertLarge / GrowSooTableToNextCapacityAndPrepareInsert on growth.
扫描 candidate = 0, 1, 2, …：探测 used 以获取 candidate 的成员身份(SwissTable 组扫描，@ 0x109d1100 … 0x109d1199)；如果不存在，则返回candidate；否则 candidate++(*v6 = v7 + 1 增量，@ 0x109d11b9 区域)并重新探测。返回第一个间隙。

text

inner_color(node, used_set):
    used = {}
    for nb in neighbors(node):           # HloPtrComparator order
        c = color_map.get(nb)
        if c is not None: used.insert(c)
    candidate = 0
    while candidate in used:             # SwissTable membership (crc32 hash)
        candidate += 1
    return candidate                     # smallest-available-color
```text

这是经典的首次拟合图形着色。颜色 0 是默认/共享屏障；由于冲突，颜色 `>= 1` 是额外的明显障碍。 The number of colors used per key equals the chromatic number of that key's interference graph (≈ the max concurrent in-flight collectives of that key), and that is exactly how many distinct barrier ids that key consumes. Because the adjacency input is a `std::map` (deterministic key order) and the search is deterministic first-fit, the coloring is **reproducible run-to-run** — no randomness — which matters for cache-stable compiles.

---

## 着色方式 `BarrierConfig` {#how-the-coloring-feeds-barrierconfig}

两个着色通道的冲突集合并为 [障碍分配生产者](overview.md) (`TensorCoreBarrierAssignment::Run`) 中的一个 `flat_hash_set<HloInstruction*>`。 In the per-key assignment loop, the entry's first instruction is looked up in that merged conflict set → `has_conflict`, the third argument to `DetermineBarrierConfigForKey @ 0x109c6fa0`:

- **color 0, NOT in conflict set** → `DetermineBarrierConfigForKey` may produce a shared `REPLICA(2)` barrier or a `GLOBAL(1)` barrier per its global-barrier-beneficial heuristic.
- **非零颜色/IN 冲突集** → 强制使用新的每键 `CUSTOM(3)` 屏障(新 id)。

This is the precise mechanism by which "two same-key TC collectives are forced apart": the coloring proves their live ranges overlap → they land in the conflict set → `DetermineBarrierConfigForKey` gives them distinct `CUSTOM(3)` `BarrierConfig`s. The `BarrierType` enum (`GLOBAL=1`/`REPLICA=2`/`CUSTOM=3`/`MEGACORE=4`) and the global-barrier semantics are detailed on the [概述](overview.md) and the [全局屏障窗口](global-barrier-window.md) pages;生成的 SFLAG *号码绑定*位于 [SFLAG 绑定](barrier-to-sflag-binding.md) 页面上。

---

## `BarrierConfig` → 芯片-SFLAG 降低

The chosen `BarrierConfig {type, id}` is written into each collective's `BackendConfig` proto by the assignment pass, then read back and lowered at kernel-emission time.决策和降低**仅**通过原型字段进行通信 - 没有共享的内存状态。

### `CustomKernelEmitter::Emit` — 读取 `BarrierConfig`，调用两个降低器

`Emit @ 0x1321ad60` (TU `custom_kernel_emitter.cc`) prepares the kernel module (`GetCustomCallConfig @ 0x13e308c0`, parse the embedded MLIR module, `SetArgLayouts @ 0x13219e80`), then:

- `HloInstruction::backend_config<BackendConfig> @ 0xf58e6c0` + `BackendConfig::CopyFrom @ 0x1d6e7400` → 提取由赋值传递写入的 `BarrierConfig` 子消息。 `r13`持有`&BarrierConfig`(如果没有则默认为`BarrierConfig_globals_ @ 0x223a9450`)； `rbp-0x548`持有`&CustomCallConfig`。
- 称为 `MaybeInsertGlobalBarrier(module, &CustomCallConfig, &BarrierConfig) @ 0x1321b2cb`。
- 调用 `RunPasses(module, b1, b2, &BarrierConfig) @ 0x1321b2f2` (两个 `b` 参数是发射模式布尔值)。

反编译确认两个调用站点和 `backend_config` + `CopyFrom` 在同一路径上读取。

### `MaybeInsertGlobalBarrier` — type-1 全局 → SFLAG 树屏障

`MaybeInsertGlobalBarrier @ 0x1321ac20` 采用 `(ModuleOp by value, CustomCallConfig const* cc, BarrierConfig const* bc)`。门(反编译字节精确)：

```c
v4 = *(int*)(cc + 0x10);                      // CustomCallConfig hasbits
// is_communicating valid iff hasbit 0x200; skip_device_barrier valid iff hasbit 0x2000
v5 = (v4 & 0x200) ? *(char*)(cc + 0x90) : 0;  // [cc+0x90] = is_communicating
v6 = (v4 & 0x2000)? *(char*)(cc + 0x94) : 0;  // [cc+0x94] = skip_device_barrier
v7 = bc ? (*(int*)(bc + 0x20) == 3) : 0;      // BarrierConfig type == 3  ⇒ fresh per-key, NOT global
// barrier-request bit = (v4 & 0x40)

三个 RET_CHECK 退出(MakeErrorImpl<3>，直接从 .rodata 读取错误字符串，确认行/偏移)：

线	VA	`.rodata` 错误字符串	触发器
`0xde8` (3560)	`@ 0x1321ac9a`	`"Custom barrier requested for non-communicating custom call."` (`@ 0xa05ba1c`)	屏障已请求 (`0x40`)，但自定义调用无法通信
`0xdf4` (3572)	`@ 0x1321acbf`	`"The compiler failed to allocate a barrier semaphore and Mosaic wasn't allowed to perform a global barrier due to skip_device_barrier."` (`@ 0xa03b568`)	每键 SFLAG 分配失败并且 `skip_device_barrier` 阻止全局回退
`0xe61` (3681)	`@ 0x1321ad48`	`"Requested barrier for unsupported core type"` (`@ 0x86b0f74`)	func-op 遍历未找到预期 `CoreType` 的操作

如果 BarrierConfig.type == 3(每个键)门提前返回 - 无全局障碍；每个键的屏障由常规 RunPasses 管道发出。否则，当需要并允许全局屏障时，落入插入路径：HasAnyCoreType @ 0x13e30700 通过 index = 2 - has_core 选择每个核心类型的屏障形状，然后 mlir::detail::walk @ 0xea26de0 over func ops 调用每个匹配 func op 的 $_0 回调 @ 0x1322a7e0：

TPUDialect::GetCoreTypeAttr @ 0x14aa6020(必须匹配，否则 WalkInterrupt →“不受支持的核心类型”RET_CHECK)。
arith::ConstantIntOp(0) / ConstantIntOp(1) @ 0x1caca3a0(计数/增量操作数)+模块配置派生的常量(GetModule()->config [+0x20])。
sparse_core::MemorySpaceAttr::get(ctx, sflag) @ 0x1458ff20 + MemRefType::get @ 0x1d897680 + arith::MulIOp @ 0x1caf0c40 (per-core SFLAG stride) → AllocateAtOffsetOp::create @ 0x145a5aa0 = the global-barrier SFLAG memref (MemorySpace::sflag).
UnrealizedConversionCastOp @ 0x1d8880e0(将 sflag memref 桥接到 TC 信号量类型)+ tpu::MemorySpaceAttr::get @ 0x14a9db60。
两个 scf::ForOp @ 0x17866d60(对等核心上的循环)包裹 tpu::SemaphoreSignalOp::create @ 0x14b442e0(在每个对等点上碰撞全局 SFLAG)，然后是 tpu::SemaphoreWaitOp::create @ 0x14b45460(等待所有对等点发出信号)。

TC 全局屏障是每个核心的 SFLAG，通过 tpu.sem_signal 向所有对等点发出信号，并通过 scf.for 循环内的 tpu.sem_wait 在核心集上等待 — 信号全部然后等待树屏障。 This is the TC analog of the SparseCore reserved global-barrier SFLAG: same SFLAG number space and the same AllocateAtOffsetOp(MemorySpace::sflag) primitive, with a different atomic op family (tpu.sem_* on TC vs sc_tpu.sync_* on SC).保留块整数位于按代号保留 SFLAG 和专用同步标志因此，页上。

LOW — 字面全局屏障槽索引。 AllocateAtOffsetOp 偏移量 Value 由 GetModule()->config 派生常量加上每核 MulIOp 步幅构建。 AllocateAtOffsetOp(MemorySpace::sflag) 原语和每核心步幅已确认，但偏移量 Value 的确切来源 — const 与 Target::GetGlobalBarrierSyncFlagNumber 访问器调用 (@ 0x1d60f420) — 在此调用站点并未隔离。文字槽被推断为保留的全局槽。

`RunPasses` — 每键(类型 2/3)降低管道

RunPasses(module, b1, b2, &BarrierConfig) @ 0x13202780 保存 BarrierConfig ([rbp-0x220] = r8 @ 0x13202794)，构建 mlir::PassManager @ 0x1cb700a0，并运行 $_0 闭合 @ 0x132048e0，该闭合安装 IR 打印 (enableIRPrinting @ 0x1cb66120，由 VLOG 门控)、PassErrorDiagnosticHandler 和通道管道：通过 OpPassManager::addPass @ 0x1cb6c000 添加 mlir::tpu::createConvertIntegerMemrefsPass @ 0x132bca60，然后通过 OpPassManager::nest @ 0x1cb6d3e0 添加嵌套通道。反编译确认了 PassManager 构造、$_0 调用和 ConvertIntegerMemrefsPass 添加。

SC 内核的每键 SFLAG 发射经过 CollectiveEmitterBase::EmitCustomBarrierFromConfig @ 0x13352cc0 / EmitCustomBarrierStart @ 0x13352fc0 → GetSyncFlagForBarrierId @ 0x133e9dc0 → AllocateAtOffsetOp(sflag) — 与 SparseCore 每环屏障使用的路径相同。因此，TC 每键屏障(BarrierConfig 类型 2/3)和 SC 每环屏障取自相同的保留芯片 - SFLAG 块； SFLAG 绑定页面记录了彩色 BarrierConfig.id 如何映射到具体的 SFLAG 编号。

低 - RunPasses 嵌套传递主体。 RunPasses 捕获 BarrierConfig 并且 SC 内核发射重用 EmitCustomBarrierFromConfig / GetSyncFlagForBarrierId 路径，但确切的嵌套传递(通过 OpPassManager::nest 添加)读取BarrierConfig 类型 2/3 并选择每个键的 SFLAG 编号(相对于全局插槽)未单独拆解。该路径上的 {lambda(Pass*,Operation*)#1} @ 0x1321cb40 是 IR 打印应该打印谓词，而不是降低障碍的传递体。
推断 — CustomCallConfig 字段名称。 bool 字段名称 [cc+0x90] = is_communicating 和 [cc+0x94] = skip_device_barrier 归因于精确控制这些字节读取的两个 RET_CHECK 错误字符串；偏移量已确认，名称未根据原始描述符字段编号进行确认。同样，CustomCallConfig hasbits 0x200/0x2000/0x40 和 [bc+0x20] 中的 BarrierConfig 类型被证明为偏移量/掩码，而不是针对原型进行字段命名。

TC vs SC：决策不同，SFLAG Sink 共享

舞台	TensorCore(本页+任务制作人)	SparseCore
去重机制	实时异步干扰图上的贪婪图着色 (`BarrierColoring::Run`)，首次拟合颜色	`FlatHashMap<SparseCoreBarrierKey, long>` 上静环键
冲突谓词	两个异步集合的实时范围重叠 (`CollectConflictInfo`)	不适用(无时间表概念)
着色	`AssignColorsGreedy @ 0x109d0c80` — 最小可用颜色	单调 `GetNextUniqueSyncFlagId`
异步开始/完成	ACP：CP-开始 `0x24` / CP-完成 `0x23`； A2A：`kBarrierStart` 自定义调用/异步启动 `0x11`	稀疏核心自定义调用
配置已制作	`BarrierConfig {type 1 global / 2 shared / 3 fresh}` → `BackendConfig`	`barrier_id` + hasbit
在发射时读回	`CustomKernelEmitter::Emit @ 0x1321ad60` → `backend_config<BackendConfig>`	`UniDirRingStrategy::BarrierId`
GLOBAL 降低	`MaybeInsertGlobalBarrier @ 0x1321ac20` → `AllocateAtOffsetOp(sflag)` + `scf.for(sem_signal/sem_wait)` 树障	`AssignGlobalBarrier` 保留槽 → 相同的 SFLAG 块
PER-KEY 降低	`RunPasses @ 0x13202780` → `EmitCustomBarrierFromConfig @ 0x13352cc0` → `GetSyncFlagForBarrierId` 型 `AllocateAtOffsetOp(sflag)`	`EmitCustomBarrierStart @ 0x13352fc0` → `GetSyncFlagForBarrierId @ 0x133e9dc0`
原子运算系列	`tpu.sem_signal` / `tpu.sem_wait`(TC 音序器)	`sc_tpu.sync_add` / `sync_wait`(SC 音序器)
SFLAG 水槽	`AllocateAtOffsetOp(MemorySpace::sflag)` → 芯片 SFLAG 块 — 双臂共享	相同

决定不同(TC = 实时冲突着色；SC = 静态环哈希)；硬件 SFLAG 接收器是相同的。 TC全局屏障使用TC树屏障(sem_signal/wait)；每个键的屏障重用 SC sync_add 发射路径。

交叉引用

概览 — 基于 SFLAG 的屏障模型和 BarrierType 枚举(1 个全局/2 个共享/3 个新鲜)。
屏障 → SFLAG 号码绑定 — 彩色 BarrierConfig {type, id} 如何映射到具体的硬件 SFLAG 编号。
InferBarrierConfig — 提供绑定的每代 SFLAG-map 源。
TensorCore 屏障 — A2A 策略开启的异步屏障操作系列(操作码 0x81 屏障启动)。
全局屏障窗口 — 此页面降低的 1 类全局屏障语义。
每个代号 compiler_reserved SFLAG — 两个屏障臂从中分配的保留芯片 SFLAG 块的文字 {base, count} 整数。
SpecialPurposeSyncFlags — 运行时 SFLAG 接收器和覆盖语义。
概述 — 障碍分配生产者 — TensorCoreBarrierAssignment::Run 运行这两个 BarrierColoring 遍，并通过 DetermineBarrierConfigForKey @ 0x109c6fa0 将冲突集转换为 BarrierConfig。
二进制： extracted/libtpu-0.0.40-cp314-cp314-manylinux_2_31_x86_64/libtpu/libtpu.so(构建 ID 89edbbe81c5b328a958fe628a9f2207d)— 返回索引。

屏障着色 ​

摘要 ​

引擎类及其两个实例化 ​

用途 ​

AssignColorsGreedy — 首次适配最小可用颜色 ​

RunPasses — 每键(类型 2/3)降低管道 ​

TC vs SC：决策不同，SFLAG Sink 共享 ​

交叉引用 ​