ICI All-Reduce 原语

本页中的所有地址、符号和偏移均适用于 libtpu-0.0.40-cp314 wheel 中的 libtpu.so（build-id 89edbbe81c5b328a958fe628a9f2207d，781,691,048 字节，完整符号，.text VMA == 文件偏移）。其他版本会不同；请将每个 VA 都视为绑定到该版本。

摘要

xla::jellyfish::AllReduceEmitter::EmitAllReduce（0x13742200，7,962 字节，1,445 行反编译代码）是 ICI fabric 上片上到片上 AllReduce 的硬件级发射入口。它不是算法，也不是选择器。它是在 SPMD partitioner 已经判定需要 AllReduce、且 SelectNDStrategy 选择器已经选定环形形状之后，构造恰好一个具体的子发射器（RingSumEmitter、UniDirection*RingStrategy，或 RotatedPincer*/AsyncPincer* 家族之一），并驱动它发射实际逐步 LLO 程序的函数：彩色环上的 reduce-scatter 然后 all-gather 循环，其中每一步是一条 remote-write DMA、一次远端 sync-flag bump、一次本地 sync-flag wait，以及一次 VPU reduce。本页负责这层分派以及逐步发射形状。

签名就是契约：

EmitAllReduce(absl::Span<const ReplicaGroup>,
              const std::function<LloValue*(LloValue*, LloValue*, LloRegionBuilder)>& fmerge)
```text

第二个参数，即**合并函子**，是 reduction 种类进入 emitter 的唯一位置。`EmitAllReduce` 从不检查它：它把这个闭包传入自己构造的任意子发射器，而子发射器在每一步对 `(local_copy, just_received_chunk)` 调用它一次。线上没有 reduction-op 字段（见 [DMA 描述符](dma-descriptor.md)）；reduce 始终是本地 VPU 操作。

熟悉 MPI 环形/递归倍增 AllReduce 的读者可以这样理解框架：这个函数把“用这些 replica groups、使用*这个*逐元素 merge 来做 reduce”转换为具体 LLO 指令流，方式是基于拓扑 + 大小 + prefer-flags 选择一个子发射器，然后遍历它的两阶段 step loop。重新实现契约如下：

- **分派分叉。** `EmitAllReduce` 首先基于 inference / size gate（`ShouldUseInferenceShortLatencyEmitter` 加 shard-size 阈值）进入 **N-D 分支**（多轴，由 `SelectNDStrategy` 驱动，反编译第 794 行）。当该 gate 不触发时，函数落入第二个分支，计算 `MayUseSinglePhaseRingEmitter`（第 997 行），然后在**二项式单阶段环**（`CreateEmitter`，第 1054 行）和 **1-D 分支**（由 `GetRingLocationWrapper` 驱动，第 1093 行）之间选择。因此二项式路径和 1-D 分支是非 N-D 分支内部的同级路径，二项式*不是*横跨 N-D 分支的快速路径。每个分支随后可选地把基础环策略包装为 pincer 或量化 pincer emitter。
- **逐步发射。** 无论选中哪个子发射器，step body 都是相同的四个线级事件：remote-write unicast DMA、远端 sync-flag bump（接收端，自动）、本地 sync-flag wait、通过 `fmerge` 执行 VPU reduce。阶段为 `Running phase 0`（reduce-scatter）、`Running phase 1`（all-gather），以及可选的 `Running phase 2`（清理）。
- **子发射器路由。** 终端类的精确集合（`RingSumEmitter`、`UniDirection1DRingStrategy`、`StrategyRing`、`UniDirectionNDRingStrategy`、`RotatedPincerEmitter`、`RotatedPincerShortEmitter`、`RotatedPincerQuantizedEmitter`）以及在它们之间路由的谓词 gate。

| | |
|---|---|
| **发射入口** | `AllReduceEmitter::EmitAllReduce` @ `0x13742200`（7,962 B / 1,445 行） |
| **顶层包装器** | `AllReduceEmitter::Emit` @ `0x13745de0`（sync → `EmitAllReduce`，fusion → `EmitAllReduceFusion`） |
| **Fusion 变体** | `AllReduceEmitter::EmitAllReduceFusion` @ `0x13746360` |
| **选择器（上游）** | `BaseStrategyND::SelectNDStrategy` @ `0x137c78e0` — 见 [strategy-nd-picker](../collectives/strategy-nd-picker.md) |
| **二项式 gate** | `RingSumEmitter::MayUseSinglePhaseRingEmitter` @ `0x1375c1c0` |
| **二项式构造器** | `RingSumEmitter::CreateEmitter` @ `0x13760720` |
| **Inference-short gate** | `RingSumEmitter::ShouldUseInferenceShortLatencyEmitter` @ `0x1375ea20` |
| **1-D 环定位器** | `AllReduceEmitter::GetRingLocationWrapper` @ `0x137412c0` |
| **跨模块测试** | `cross_replica_sharding_util::IsCrossModuleReduceInstruction` |
| **支持的 dtype** | `kSupportedTypes` @ `.rodata 0x0ae5a56c` = `{F32=11, S32=4, U32=8, BF16=16, PRED=1}` |
| **源 TU** | `platforms/xla/service/jellyfish/lowering/`（`ring_sum_emitter.cc`、`rotated_pincer_emitter*.cc`） |
| **置信度** | 对分派分叉和子发射器集合为 HIGH（调用点经反编译验证）；逐行例外已注明 |

---

## 所处位置

`EmitAllReduce` 是 **AllReduce 编译栈的底部**，位于各家族线级 emitter 之上一层。链路如下：

```text
HloAllReduce  ──(SPMD partitioner, replica groups decided)──▶
  AllReduceEmitter::Emit  (0x13745de0)         # sync vs fusion fork
    ├─ sync     ──▶ AllReduceEmitter::EmitAllReduce       (0x13742200)   ◀── THIS PAGE
    └─ fusion   ──▶ AllReduceEmitter::EmitAllReduceFusion (0x13746360)
                          │
   (inside EmitAllReduce)  ├─ SelectNDStrategy / GetRingLocationWrapper  # choose ring shape
                          ├─ construct ONE sub-emitter (ring / pincer)
                          └─ drive its 2-phase step loop (RS + AG)
                                 per step: DMA → remote sflag bump → local sflag wait → VPU reduce
                                              │
                                  EnqueueDmaInGranules → DMA_TYPE_REMOTE_WRITE_UNICAST  (../ici/dma-descriptor.md)

本页不覆盖的内容及链接：

collective 级二项式 / 递归倍增算法，即逐 rank butterfly partner 调度、int32[N×8] replica 表、可行性 gate，位于二项式 / 递归倍增。本页只记录 EmitAllReduce 会路由到它，以及如何路由。
策略选择器，即把拓扑分类为五种终端环类的决策树，是 SelectNDStrategy，记录于 SelectNDStrategy。EmitAllReduce 调用它；并不重新实现它。
双向 pincer 循环形状（重叠的 send/recv 窗口、[dim][color] sflag 表）位于分层 AllReduce / Pincer。
每代 DMA 描述符字节布局以及远端 sync-flag 编码位于 DMA 描述符。
量化 8-bit 线级路径（RotatedPincerQuantizedEmitter、{S8, F8E5M2, F8E4M3B11FNUZ} 线级集合）位于 FP8 量化 Collective；本页只注明进入它的分派 hook。

分派分叉

EmitAllReduce 是一个很高的 if/else 级联。它验证 instruction 并读取其 BackendConfig（第 656 行的 HloInstruction::backend_config<jellyfish::BackendConfig>，携带 BarrierConfig 以及 partition 时烘入的任何 prefer-flags）之后，把世界切成三条互斥路径。下面的反编译位置编号来自 0x13742200。

顶层分支是第 775 行的 N-D / 非 N-D 分裂（if (!ShouldUseInferenceShortLatencyEmitter && !(size_gate | …))）。N-D 分支位于taken 侧（下面 Step 2，第 794 行的 SelectNDStrategy）；二项式和 1-D 分支是 else 侧（第 985 行以后）中的同级路径，由第 1010 行的 MayUseSinglePhaseRingEmitter 测试分隔。下面的 Steps 按二项式 → N-D → 1-D 排列是为了叙述清晰，不是源码顺序；实际控制流请查看行号。

Step 0 — 序言：replica groups、跨模块测试、barrier config

text

emit_all_reduce(replica_groups, fmerge):                     # 0x13742200
    is_cross_module = IsCrossModuleReduceInstruction(hlo)     # line 557 — replica × partition fold
    cfg             = hlo.backend_config<BackendConfig>()      # line 656
    barrier         = cfg.has_barrier ? cfg.barrier            # line 682/688 — BarrierConfig
                                      : BarrierConfig_globals_  #   default global barrier
    # (the merge functor `fmerge` is captured here, never inspected)
```text

`IsCrossModuleReduceInstruction`（第 557、560 行）把 AllReduce 分类为跨模块（同时跨 replicas *和* partitions reduce）或单模块。它作为布尔值 `IsCrossModuleReduceInstruction`（存于 `[rbp-2Ah]`）被读取，并传入每个下游构造器；它会改变环定位器使用的 replica-group fold，并 gate 二项式跨模块子路径。

### Step 1 — 二项式单阶段环（1-D 分支的同级，第 1010 行以后）

```text
    may_single = RingSumEmitter::MayUseSinglePhaseRingEmitter(mem_unit,     # 0x1375c1c0, line 997
                                  span_size, max_scratch, target, env)
    if env.force_1d_ring != 1 or (size_gate | short_latency                # line 1010 — branch test
                                  | may_single | (num_colors != 3)):
        # build the binomial / single-phase ring emitter; install two closures:
        ring_loc_provider = $_0  # returns net_util::RingLocation   (line 1031)
        group_provider    = $_1  # returns BinomialGroupData         (line 1039)
        emitter = RingSumEmitter::CreateEmitter(span_size, max_scratch,      # 0x13760720, line 1054
                                                ..., fmerge, barrier, ...)
        emitter->Build(); emitter->Emit()                                    # virtual, vtable+8 / vtable+...
        return

MayUseSinglePhaseRingEmitter（0x1375c1c0）是第 1010 行分支测试中的一项；除非 env 的 force-1D-ring 标志已设置且 {size_gate, short_latency, may_single, num_colors!=3} 均不成立，否则会进入二项式/环侧；若这些条件都不成立，控制流会落入 1-D 分支。CreateEmitter（0x13760720）随后构造具体的 BinomialSinglePhaseRingSumEmitter（当可行性标志已设置时）或普通 SinglePhaseRingSumEmitter 环（两个 SetEmitter tag，即 "BinomialSinglePhaseRingSumEmitter" 和 "SinglePhaseRingSumEmitter"，都存在于 CreateEmitter 函数体中）。两个闭包 $_0 和 $_1 是二项式数据路径的关键证据：$_0 产生 net_util::RingLocation（core 的环位置），$_1 产生 BinomialGroupData（从 int32[N×8] replica 表读取的预计算 counterparts vector）。这些闭包供给的算法，即递归倍增 butterfly，完整记录在二项式 / 递归倍增；本页唯一主张是 EmitAllReduce 会构造该 emitter 并提供这两个 provider。

注意 — MayUseSinglePhaseRingEmitter 是 prefer gate，不是正确性 gate。可行性约束（N 为 2 的幂，N ≤ 128）在二项式 emitter 内部强制执行，而不是这里。如果 gate 通过但环不是 2 的幂，CreateEmitter 会回退到环形 RingSumEmitter。重新实现者不得把二项式路径理解为“小尺寸时总是采用”。

Step 2 — N-D 分支（多轴 torus）

当二项式快速路径未采用且拓扑为多轴时，EmitAllReduce 调用选择器，然后基于结果构造一个环或 pincer。

text

    strategy = SelectNDStrategy(target, env, !is_cross_module, hlo, b, ...)  # 0x137c78e0, line 794
    if want_pincer:                                                          # prefer-flag driven
        if quantized:  emitter = make_unique<RotatedPincerQuantizedEmitter>(strategy, ...,  # line 886
                                       fmerge, barrier, primitive_type, QuantizedAllReduceStage, ...)
        else:          emitter = make_unique<RotatedPincerEmitter>(strategy, ...,            # line 907
                                       fmerge, barrier, ...)
        # log tag "Long Pincer" (line 871) or "2-D rotated pincer" (line 630)
    elif short_latency_pincer:
        emitter = RotatedPincerShortEmitter::CreateIfFeasible(strategy, ..., barrier, ...)   # line 847
    else:
        strategy_obj = new UniDirectionNDRingStrategy(strategy, b)   # operator new(0x5B0), line 815/816
        emitter      = strategy_obj                                  # the ring IS the emitter
    emitter->Build()                                                 # (*(vtable+8))(emitter), line 921
    emitter->Emit()
```text

`SelectNDStrategy` 返回 `BaseStrategyND*`（反编译中的 `v93`/`v136`）。随后 N-D 分支要么：

- 通过把该 strategy 对象包装成 `UniDirectionNDRingStrategy`（第 815 行的 `operator new(0x5B0)`，即 1,456 字节的 ND-ring emitter 对象）来直接使用它，或
- 把 strategy 交给 pincer 构造器（第 907 行的 `RotatedPincerEmitter`）或其量化变体（第 886 行的 `RotatedPincerQuantizedEmitter`，还接受 `PrimitiveType` 和 `QuantizedAllReduceStage`），或
- 通过 `RotatedPincerShortEmitter::CreateIfFeasible`（第 847 行）交给 short-latency pincer。

这些选择由 prefer-flag 驱动（选择器读取的也是这些标志）；`EmitAllReduce` 只读取结果。无论构造的是哪个对象，都会通过相同的虚拟 `Build()`/`Emit()` 对驱动（第 921 行的 `(*(void (**)(...))(*(_QWORD*)v136 + 8LL))(v136)` 间接调用即 `vtable[1]`）。

### Step 3 — 1-D 分支（单轴环）

当拓扑解析为单个可用轴时，`EmitAllReduce` 进入由 `GetRingLocationWrapper` 锚定的 1-D 路径。

```text
    ring_loc = GetRingLocationWrapper(hlo, replica_groups, is_cross_module,  # 0x137412c0, line 1093
                                      env, ...)                              #   → net_util::RingLocation
    if want_unidir:
        emitter = make_unique<UniDirection1DRingStrategy>(ring_loc, span,    # line 1120
                                                          b, shared_registry, is_bf16)
    else:
        emitter = make_unique<StrategyRing>(ring_loc, span, shared_registry, # line 1134
                                            local_value, is_bf16)
    # optional wrap, same as the N-D arm:
    if short_latency:  emitter = RotatedPincerShortEmitter::CreateIfFeasible(...)   # line 1170
    if quantized:      emitter = make_unique<RotatedPincerQuantizedEmitter>(...)    # line 1227
    elif want_pincer:  emitter = make_unique<RotatedPincerEmitter>(...)             # line 1241
    emitter->Build(); emitter->Emit()                                # (*(vtable+8))(v93), line 1205

GetRingLocationWrapper（0x137412c0）包装 GetRingLocation（0x13740e40）和 GetRingLocationWithReordering（0x137410e0），返回当前 core 在其 color 环中的 net_util::RingLocation{ring_index, position_in_ring, ring_size}，并在 routing table 需要时应用 limited-ICI-routing reorder（见路由概览）。基础环策略随后要么是 UniDirection1DRingStrategy（固定 +1 CW 邻居，第 1120 行），要么是 StrategyRing（第 1134 行），并带有与 N-D 分支相同的可选 pincer / 量化包装。

注意 — 1-D 和 N-D 分支共享相同的终端 pincer 类（RotatedPincerEmitter、RotatedPincerShortEmitter、RotatedPincerQuantizedEmitter）。区别只在它们包装的基础 strategy 对象：由 RingLocation 构建的 UniDirection1DRingStrategy/StrategyRing（1-D）对比由 SelectNDStrategy 构建的 UniDirectionNDRingStrategy（N-D）。pincer 是环上的 decorator，不是单独的环算法。

子发射器路由表

EmitAllReduce 能构造的终端子发射器全集，以及反编译调用点和 gate。所有 make_unique/operator new/CreateIfFeasible 位置均确认存在于 0x13742200。

子发射器	构建方式	分支	Gate
`BinomialSinglePhaseRingSumEmitter`	`RingSumEmitter::CreateEmitter` @ line 1054	二项式	`MayUseSinglePhaseRingEmitter` + 2 的幂 `N≤128`
`SinglePhaseRingSumEmitter`（普通环）	`RingSumEmitter::CreateEmitter` @ line 1054	二项式	`MayUseSinglePhaseRingEmitter`，非二项式回退
`UniDirection1DRingStrategy`	`make_unique` @ line 1120	1-D	单轴，单向
`StrategyRing`	`make_unique` @ line 1134	1-D	单轴，环（非 unidir）
`UniDirectionNDRingStrategy`	`operator new(0x5B0)` @ line 815	N-D	多轴，无 pincer
`RotatedPincerEmitter`	`make_unique` @ lines 907 / 1241	两者	带宽受限，pincer prefer-flag
`RotatedPincerShortEmitter`	`CreateIfFeasible` @ lines 847 / 1170	两者	延迟受限双向
`RotatedPincerQuantizedEmitter`	`make_unique` @ lines 886 / 1227	两者	`CanLowerToQuantizedAllReduce`（第 699 行）

RotatedPincerQuantizedEmitter::CanLowerToQuantizedAllReduce 会被较早查询（第 699 行）；如果量化级别和大小阈值满足且 dtype 位于 kSupportedQuantizationTypes（{S8, F8E5M2, F8E4M3B11FNUZ}）中，无论进入哪个分支都会采用量化包装。构造器比非量化 pincer 多接受一个 PrimitiveType 和一个 QuantizedAllReduceStage 参数（比较第 886 行与第 907 行的 mangled make_unique 签名）。详见 FP8 量化 Collective。

注意 — RotatedPincerQuantizedEmitter/RotatedPincerEmitter 构造器按值接受 RotatedPincerEmitterBase::ExtraArgs（可从 mangled name ...RotatedPincerEmitterBase9ExtraArgsE 看出）。其中打包了每个 color 的 VMEM scratch 大小和保留的 sflag block，使 pincer 能确定其 [dim][color] recv-flag 表大小，即双向性 bookkeeping。async-pincer 家族（AsyncPincerEmitter）在 fusion 路径（EmitAllReduceFusion，0x13746360）中构造，而不在 sync EmitAllReduce 函数体中；sync 函数体最多到 rotated pincer。

逐步发射

无论 EmitAllReduce 构建哪个子发射器，它发射的 LLO 程序都具有相同形状：一个两阶段循环，其中每一步是跨一条 ICI 链路的一次往返加一次本地 reduce。阶段字符串 Running phase 0、Running phase 1、以及（某些 emitter 中的）Running phase 2，再加上 Resetting reduction buffer for phase 0，是这些阶段的二进制锚点。

text

# the shape the chosen sub-emitter's Build()/Emit() produces
for color in 0 .. num_colors-1:                      # concurrent rings, one per torus axis
    N   = ring_size[color]
    pos = ring_position[color]                        # from RingLocation / BinomialGroupData

    # ---- (optional) startup barrier ----
    if barrier.scope != none:                         # BarrierConfig from backend_config
        BarrierStart(barrier_sync_flag, ...)          # TreeBarrierType::kAll / kCrossReplica

    # ---- PHASE 0: reduce-scatter (N-1 steps, ring) / log2(N) steps (binomial) ----
    for i in 0 .. steps-1:
        peer = next_peer(pos, i, color)               # CwCore (ring) or counterparts[i] (binomial)
        EnqueueDmaInGranules(out_shard, peer,         # DMA_TYPE_REMOTE_WRITE_UNICAST
                             remote_sflag_handle)      #   carries the receiver's sflag to bump
        recv = DmaDoneInGranules(...)                  # local sync-flag WAIT ("shard-{cw,ccw}-recv-wait")
        fmerge(local_copy, recv, b)                    # VPU vadd/vmin/vmax/vmul/vand/vor

    # ---- PHASE 1: all-gather (N-1 steps; absent on binomial, which self-completes) ----
    for i in 0 .. N-2:
        peer = next_peer(pos, i, color)
        EnqueueDmaInGranules(reduced_shard, peer, ...)
        DmaDoneInGranules(...)                          # "shard-send-wait"
        SafeMemcopyN(recv -> shard)                     # plain copy, NO merge
```text

每个 reduce-scatter step 的四个线级事件为：

1. **Remote-write DMA** — `EnqueueDmaInGranules` 发射一个 `DMA_TYPE_REMOTE_WRITE_UNICAST` 描述符：source = 出站 shard 的本地 VMEM offset，destination = 邻居 reduce-scatter scratch 中的槽位，外加一个**远端 sync-flag handle**。描述符 word 布局随代际变化；见 [DMA 描述符](dma-descriptor.md)。
2. **远端 sync-flag bump** — 当 chunk 到达时，接收方 NodeFabric Ingress Unit 会自动递增已编码的远端 sync flag（线级 `atomic_remote_add_set_done`）。编译器不会在接收端发射软件 ack。
3. **本地 sync-flag wait** — `DmaDoneInGranules` 发射对本地接收 sync flag 的 wait，注解为 `shard-cw-recv-wait` / `shard-ccw-recv-wait`（reduce 侧，方向由环 leg 限定）或 `shard-send-wait` / `shard-cw-send-wait` / `shard-ccw-send-wait`（gather 侧）。运行时 watchdog 字符串为 `Sflag wait timeout on op {…}`。
4. **VPU reduce** — wait 清除后，合并函子 `fmerge` 在 `(local_copy, just_received_chunk)` 上被调用。这是普通标量 VPU 操作（`vadd`/`vmin`/`vmax`/`vmul`/`vand`/`vor`），折叠在 step body 中。线上没有 reduce。

all-gather 阶段用普通 `SafeMemcopyN` 替代 `fmerge`，即让已经 reduce 完成的 shards 沿环旋转而不再 reduce。

> **怪癖 —** 合并函子 `fmerge` 是 `EmitAllReduce` 的*第二个参数*，并被不透明捕获。该函数从不基于 reduction kind 分支；kind 已在上游（`xla::MakeReductionComputation(kind, dtype)`）解析进闭包函数体。这是 ICI fabric 与 reduction-op 无关的结构性原因：等 `EmitAllReduce` 运行时，“SUM vs MAX”已经是 `fmerge` 内部编译好的 VPU 指令，而不是 emitter 或线上能看见的值。

---

## 输入验证和 Dtype Gate

`EmitAllReduce` 在分派之前拒绝不支持的 element type。`kSupportedTypes`（`.rodata 0x0ae5a56c`，20 字节 = 五个 `int32`）正是 `{F32=11, S32=4, U32=8, BF16=16, PRED=1}`（XLA `PrimitiveType` enum 值，字节确认）。其他任何类型都必须在上游提升，SPMD partitioner 会插入 convert。

布尔（`PRED`）路径进一步受 MLIR 级 verifier 限制：字符串 `Vector mask all-reduce must have i32 output` 和 `Mask all-reduce only supports sum and find_first_set kinds`（从 `mlir::tpu::AllReduceOp::verify` @ `0x14b01460` 恢复）把 mask AllReduce 限制为 SUM 和 find-first-set reduction。五个 `REDUCTION_KIND_*` enum 值（`UNSPECIFIED=0, SUM=1, PRODUCT=2, MIN=3, MAX=4`）是 merge functor 可编码的 kind；`UNSPECIFIED` 在验证阶段被拒绝。

BF16 accumulation 选择，即 BF16 是在线上原生 reduce 还是升级到 F32，由上游决定（`bf16_inside_cross_replica_sum` @ `0x1373ca60`，加 SPMD 标志 `xla_tpu_spmd_f32_accum_for_bf16_ar`），所以等 `EmitAllReduce` 运行时，它看到的 dtype 就是环承载的 dtype。reduce 使用的逐步 VPU mnemonic（`VADDBF16rr_V0` vs `VADDrr_V0` 等）是 `fmerge` 的属性，而不是这个函数的属性。

---

## Sync vs Fusion 分叉（`Emit` @ `0x13745de0`）

公开入口是 `AllReduceEmitter::Emit`（`0x13745de0`，1,112 字节），这是一个薄包装器，会解析 instruction 的 operands，并基于 AllReduce 是否 fused 分叉：

```text
Emit():                                                   # 0x13745de0
    parse operands / shape
    if not fused:  EmitAllReduce(replica_groups, fmerge)   # 0x13742200, line 139 — THIS PAGE
    else:          EmitAllReduceFusion()                   # 0x13746360, line 144

反编译精确显示了两个终端调用：EmitAllReduce（第 139 行）和 EmitAllReduceFusion（第 144 行）。fusion 变体（EmitAllReduceFusion @ 0x13746360，7,136 字节）处理 async（kAsyncStart/kAsyncUpdate/kAsyncDone）以及 continuation-fusion 形式，也是 AsyncPincerEmitter 家族和 EmitColorwiseFusedAllReduce（0x1374c140）所在之处。这里记录的 sync EmitAllReduce 是单个阻塞调用点：所有 reduce-scatter + all-gather steps 都在一个 LLO region 内发射。

注意 — Emit 两个分支之间的区别是 AllReduce 的 async-ness，由 SPMD TpuAsyncCollectiveCreator pass 决定。非 async AllReduce 总是进入 EmitAllReduce；async AllReduce 进入 EmitAllReduceFusion。两个分支最终都会驱动相同的子发射器家族和相同的逐步线级格式，只是 fusion 分支会把 steps 与周围 compute 交错，并通过 sflags 同步，而不是阻塞。

重新实现检查清单

将合并函子 fmerge 作为不透明 std::function<LloValue*(LloValue*, LloValue*, LloRegionBuilder)> 接收；绝不要在 emitter 内部基于 reduction kind 分支，因为它已经编译进闭包。
在分派之前拒绝 {F32, S32, U32, BF16, PRED} 之外的任何 element type；将 PRED 路径限制为 SUM / find-first-set。
在选择子发射器之前从 IsCrossModuleReduceInstruction 读取 is_cross_module，并从 instruction 的 backend config 读取 BarrierConfig；两者都会供给构造器。
按顺序通过三个检查分派：(1) MayUseSinglePhaseRingEmitter → CreateEmitter（二项式/环快速路径，安装 RingLocation 和 BinomialGroupData provider）；(2) 多轴拓扑 → SelectNDStrategy → UniDirectionNDRingStrategy（可选 pincer 包装）；(3) 单轴 → GetRingLocationWrapper → UniDirection1DRingStrategy/StrategyRing（可选 pincer 包装）。
在 1-D 和 N-D 分支中，以完全相同的方式把 pincer / quantized-pincer 包装作为所选基础环策略之上的 decorator 应用。量化包装之前先查询 CanLowerToQuantizedAllReduce。
对每个构造出的子发射器，都通过相同的虚拟 Build()/Emit() 对（vtable+8）驱动。逐步 body 必须在每个 reduce-scatter step 中发射：EnqueueDmaInGranules(REMOTE_WRITE_UNICAST) → DmaDoneInGranules（shard-{cw,ccw}-recv-wait）→ fmerge；all-gather step 用 SafeMemcopyN 替代 fmerge。
async / continuation-fusion 形式属于 EmitAllReduceFusion，不属于这里；EmitAllReduce 只是单 region 阻塞路径。

交叉引用

ICI 概览 — AllReduce 原语在 ICI 子系统中的位置（bring-up → discovery → transfer）。
DMA 描述符 — 每一步发射的各代描述符 word 布局和远端 sync-flag 编码。
链路启动 — 任何 DMA 可以承载之前将链路拉起的 firmware PHY / host DL 状态机。
故障恢复 — Sflag wait timeout watchdog、VerifyDmaCountDone，以及阻塞 AllReduce 背后的 link-fatal 级联。
二项式 / 递归倍增 — 本页路由进入的二项式快速路径背后的 butterfly partner 调度和 replica 表。
SelectNDStrategy — 把拓扑分类为 EmitAllReduce 所构造环形形状的上游选择器。
分层 AllReduce / Pincer — RotatedPincer* 包装生成的双向 pincer 循环。
Reduce-Scatter — 作为独立 collective 的两阶段分解中的 reduce-scatter 半段。
FP8 量化 Collective — 量化包装分派进入的 RotatedPincerQuantizedEmitter 8-bit 线级路径。
路由概览 — GetRingLocationWrapper 的 peer 如何变为 ICI 可路由的远端地址（limited-routing reorder）。
返回索引

ICI All-Reduce 原语 ​

摘要 ​

分派分叉 ​

Step 0 — 序言：replica groups、跨模块测试、barrier config ​

Step 2 — N-D 分支（多轴 torus） ​

子发射器路由表 ​

逐步发射 ​

重新实现检查清单 ​

交叉引用 ​

ICI All-Reduce 原语

摘要

分派分叉

Step 0 — 序言：replica groups、跨模块测试、barrier config

Step 2 — N-D 分支（多轴 torus）

子发射器路由表

逐步发射

重新实现检查清单

交叉引用