SampleCombiner 发射器

本页上的每个 emitter 地址、op 创建序列、源码行标签和操作数身份断言，都从 libtpu-0.0.40-cp314 wheel 中的 libtpu.so 按字节精确读取（build-id 89edbbe81c5b328a958fe628a9f2207d，build libtpu_lts_20260413_b_RC00）——来源包括 SparseDenseMatmulDotCombinerEmitter::{Emit, EmitSampleCombiner, EmitValencyLoop, EmitVectorizedLoop} 的反编译 body，以及每个 body 内嵌的 GetCurrentLocation 源码行字符串。.text VMA 等于其文件偏移（0xe63c000）。地址适用于此 build；其他版本会有所不同。

摘要

SampleCombiner emitter 是 SparseDenseMatmulDotCombinerEmitter 的内层循环体，也就是把一个 embedding-lookup sample（由可变长度 table-row id 列表和逐 id gain 组成）lower 成一个 dense output vector 的逻辑。它是 embedding datapath 的 reduce 阶段：stream gather/scatter 在 HBM 和 TILE_SPMEM 之间移动 row，VEX 是跨 lane scan engine，而这个 emitter 是逐 sample 的 gather-multiply-accumulate（GMA）循环，把一个 sample 的 rows 合并成它的 activation row。决定性结构事实是，三个 emitter 形成一个严格的三层 loop nest：EmitSampleCombiner ⊃ EmitValencyLoop ⊃ EmitVectorizedLoop，而不是运行时选择的一组替代路径。它们之间没有分支；每一层都无条件调用下一层。

第二个决定性事实是，combiner 名册收缩成一条代码路径。此 emitter 执行的 reduction 总是 weighted sum：对 sample 的 CSR segment 中每个 id，gather 它的 embedding row，把该 id 的 gain broadcast 成 feature-width vector，并把 row · gain fused-multiply-add 进清零的 SPMEM accumulator。三个具名 pooling combiner：sum、mean、sqrtn，不是不同代码路径，也不是 backend-config 字段。它们只在 front-end（XLA custom-call 之上的 TF/JAX TPU-embedding layer）折叠进 sorted_gains 操作数的逐 id gain 值上不同：sum 保持 gain 不变，mean 预先按 1/valency 缩放，sqrtn 按 1/sqrt(valency) 缩放。emitter 会原样把交给它的 gain 当作乘法 scale。通用（非 pooling）combiner 位于独立的 CustomCombiner family 中，它会以内联 user reduce computation 替换固定 FMA。

本页负责三件事：combiner reduction 名册（固定 weighted-sum FMA、gain-as-divisor 约定，以及 DotCombiner vs CustomCombiner vs 非 minibatch GatherMulScatter family 拆分）、逐 combiner weight 应用（UnalignedLoadScalarFromHbm → BitcastOp i32→f32 gain load、BroadcastScalarToVector fan-out，以及 MulFOp/AddFOp FMA），以及内层循环发射（带 bounds-guard 的 Emit sample-tile loop、EmitSampleCombiner accumulator scope，以及 EmitValencyLoop 逐 id loop）。它消费的 CSR multiplicity / valency 由 Dedup Multiplicity 和 EmitValencyLoop 负责；它发出的同步 gather 由 stream gather/scatter 负责。它们在此链接，不重复说明。

对重新实现来说，契约是：

combiner 是固定 weighted-sum FMA。 acc_chunk += emb_chunk · broadcast(gain)，FastMathFlags = none。sum/mean/sqrtn 是 gain scale，不是代码路径，也不是配置字段。重新实现者只发射一个 FMA loop，并信任 front-end 已经预先缩放 gain。
三个 emitter 是严格嵌套，不是 dispatch。 EmitSampleCombiner（逐 sample accumulator）→ EmitValencyLoop（逐 id loop）→ EmitVectorizedLoop（逐 chunk FMA）。运行时不会在“valency loop”和“vectorized loop”之间选择；vectorized loop 是 valency loop 的 body。
gain 按原始 i32 bits 加载并 bit-cast 为 f32。 逐 id gain 由 UnalignedLoadScalarFromHbm 获取（返回整数标量），再通过 arith::BitcastOp 重新解释为 f32。它绝不是 arith::SIToFPOp 风格的数值转换；HBM word 就是 float 的 bit pattern。
accumulator 是清零的 SPMEM tile，作用域由 memref::AllocaScopeOp 限定。 EmitSampleCombiner 分配一个 f32 scoped buffer，通过 InitializeTileSpmemBuffer（ZeroMemOp）清零，用该 buffer 作为 iter_arg 运行 valency loop，然后用同步 stream op 把它 drain 到 HBM。
Emit 的 scf::IfOp 是 tile bounds guard，不是 minibatch dispatch。 它的 then region 运行 EmitSampleCombiner；else region 为空。它跳过最后一个部分 sample tile 中越界的 lane；它不是 num_minibatches == 1 分支。


Emitter class	`xla::tpu::sparse_core::SparseDenseMatmulDotCombinerEmitter`
Source file	`platforms/xla/sparse_core/sparse_dense_matmul_dot_combiner_emitter.cc`（来自 `GetCurrentLocation` 字符串）
Entry / nest	`Emit` `0x1332bda0` ⊃ `EmitSampleCombiner` `0x1332c640` ⊃ `EmitValencyLoop` `0x1332cee0` ⊃ `EmitVectorizedLoop` `0x1332e1c0`
Dispatcher	`LoweringEmitter::EmitSparseDenseMatmulDotCombiner` `0x131a7ca0`（构建 2 个 `FoldAllDimensions` 值 → ctor → `Emit`）
Combiner reduction	固定 weighted-sum FMA `acc += emb · gain`；`MulFOp` + `AddFOp`，`FastMathFlags = none`
Combiner roster	`sum` / `mean` / `sqrtn` = 一条路径，由 front-end gain scale 区分（不是代码路径，不是配置字段）
Weight source	`sorted_gains` 操作数；逐 id scalar 通过 `UnalignedLoadScalarFromHbm` → `arith::BitcastOp` i32→f32
Accumulator	f32 SPMEM tile，`memref::AllocaScopeOp` + `AllocateScopedMemory` + `InitializeTileSpmemBuffer`（`ZeroMemOp`）
Gather	`InitiateSynchronousStreamOperation`（= `LinearStreamStartOp` + `StreamWait` + `SetSyncFlag`），由逐 id token offset keyed
置信度	CONFIRMED（反编译 op-sequence 和源码行锚定），除非某行或 callout 另有说明

NOTE — 本页负责 combiner reduction、weight application 和 inner-loop op sequence。 该 loop 迭代的 CSR multiplicity / valency 由 Dedup Multiplicity 负责；孤立的逐 id loop 结构由 EmitValencyLoop 负责；同步 indirect gather / scatter-add primitive 由 stream gather/scatter 负责；此 lowering 之上的 HLO minibatching decomposition 由 Embedding Minibatching 负责。它们在此链接，不重复。

Combiner Reduction 名册

一个 reduction，三个名称

DotCombiner emitter 执行单一算术 reduction：对 gather 到的 embedding rows 做按 gain 加权的 sum。在反编译的 EmitVectorizedLoop body 中，核心精确地是两个 arith op，顺序如下：

text

EmitVectorizedLoop core (0x1332e1c0, src lines 250 / 252)
  mul = MulFOp(emb_chunk, broadcast_gain)   ; src ln 250 — emb · gain
  acc = AddFOp(mul, accumulator_chunk)       ; src ln 252 — + running accumulator
        StoreChunk(accumulator, acc)         ; no .cc line tag — inherits AddFOp location
```text

`MulFOp::create` 和 `AddFOp::create` 都以 `FastMathFlags = none` 发射（调用前 flags byte register 被 xor 清零）。这个 body 里没有除法、没有平方根，也没有条件分支；reduction 无条件是 `acc += emb · gain`。

TPU-embedding feature 可请求的三个 pooling combiner 会映射到这一个 FMA，如下所示。combiner divisor 在到达 XLA custom-call 之前就**折叠进逐 id gain**；emitter 从不看到 combiner 名称：

| Combiner | front-end 提供的逐 id gain | emitter 执行的操作 |
|---|---|---|
| `sum` | `gain` unchanged | `acc += emb · gain` |
| `mean` | `gain · (1 / valency)` | `acc += emb · gain`（相同 op） |
| `sqrtn` | `gain · (1 / sqrt(valency))` | `acc += emb · gain`（相同 op） |

> **NOTE — `sum`/`mean`/`sqrtn` 是 gain scale，不是代码路径。** `SparseDenseMatmulConfig` backend-config 中没有 `combiner_type` 字段，emitter 也没有基于 combiner enum 的分支。divisor（`1/n`、`1/sqrt(n)`）应用在 XLA layer 之上，由 TF/JAX TPU-embedding API 通过预先缩放每个 id 的 gain 完成。计算 `gain = 1/n` 或 `1/sqrt(n)` 的 front-end op **不**存在于 `libtpu.so`（它位于 custom-call 之上的 layer）。二进制中 CONFIRMED 的事实是 emitter 原样乘以 gain；该 gain *携带* combiner divisor 是从缺少任何 combiner 字段和 API 约定 INFERRED 而来。

### Combiner emitter family

`DotCombiner` 是二进制中携带的三种 embedding-combine lowering 之一。每一种都是不同的 emitter / decomposer：

| Emitter / decomposer | Address | Reduction form |
|---|---|---|
| `SparseDenseMatmulDotCombinerEmitter::Emit` | `0x1332bda0` | fixed weighted-sum FMA（本页） |
| `LoweringEmitter::EmitSparseDenseMatmulDotCombiner` | `0x131a7ca0` | dispatcher：2 个 `FoldAllDimensions` operands → ctor → `Emit` |
| `XlaSparseDenseMatmulCustomCombinerOnTc{,Grad}WithCsrInputOp`（+ `…GradWith{Sgd,Adagrad,AdagradMomentum,…}AndCsrInputOp`） | `0xe653640`（Compile）ff. | 以内联 user reduce computation 替代 FMA |
| `GatherMulScatterSparseDenseMatmulOpDecomposer`（`AddPass`） | `0x1306d740` | non-minibatch HLO decomposition（无 CSR-window decomposition） |
| `SparseGatherEmitter::Emit` | `0x133f9120` | standalone SC gather（non-minibatch 路径 lower 经由的 gather half） |

`CustomCombiner` family 是泛化版本：它不是固定的 `acc += emb · gain`，而是以内联 user-supplied `HloComputation` reduce 处理 gather 到的 rows（它携带用于 forward/backward reduce 的 `BuildCombinerVjpComputation` 和 `EmitSparseCoreComputations` 成员）。`GatherMulScatter` decomposer 是当 CSR minibatching decomposition *未*应用时 embedding lower 到的更简单对应路径：同一个 sum-lookup 的另一种 lowering，在此二进制中已命名但本页没有 body-decode。

> **WARNING — `DotCombiner` 路径不使用 Sort/Unique/SegmentedScan dialect DAG。** 一个合理假设是 embedding sum-lookup 总是通过跨 lane [scan datapath](scan-datapath.md)（`SegmentedScanOp` 等）lower。这个 emitter **不是**这样：它使用 scalar/streaming combiner，即逐 id `UnalignedLoadScalarFromHbm` + 同步 `LinearStream` gather + 对 SPMEM accumulator 的 `arith` FMA。segmented-scan dialect 路径是 embedding reduce 的*另一种* lowering；`DotCombiner` 是 gather-FMA 形式。

---

## 逐 Combiner Weight 应用

### 加载 gain

combiner weight 是 `SparseDenseMatmul` op 的 `sorted_gains` 操作数：每个（sample, id）对一个 f32 gain，其布局与 CSR segment 索引的 sorted-id list 平行。在 `EmitValencyLoop` 中，逐 id gain 通过两个 op 获取并转换：

```text
EmitValencyLoop gain load (0x1332cee0, inside the per-id scf::ForOp body)
  g_i32 = UnalignedLoadScalarFromHbm(gains_base, token_offset)   ; raw 32-bit word (no .cc line tag)
  g_f32 = arith::BitcastOp(f32, g_i32)                           ; getF32Type + BitcastOp — reinterpret bits (LocationGenerator variant loc, no integer line)

该 load 返回一个整数标量；bit pattern 通过 arith::BitcastOp 重新解释为 f32（反编译中 mlir::Builder::getF32Type 紧邻在 BitcastOp::create 之前）。这是bit reinterpretation，不是数值转换；HBM word 直接保存 gain 的 IEEE-754 bits。如果重新实现者在这里发射整数到浮点的数值转换，会产生错误 activation。

Broadcast 并应用 gain

scalar gain 会在每个 id 上、chunk loop 外部 broadcast 成 feature-width vector，然后乘入每个 chunk：

text

EmitVectorizedLoop weight application (0x1332e1c0)
  ConstantIndexOp(chunk_count)                         ; src ln 223 / 225 / 226 — loop bounds setup
  bcast = BroadcastScalarToVector(lane_width, g_f32)   ; gain → vector, hoisted out of the loop
  for chunk in [0, feature_width):                     ; scf::ForOp (src ln 234), iter_arg = accumulator chunk
    emb   = LoadChunk(gathered_row,  chunk)            ; the gathered embedding chunk (operand v3 base)
    a_cur = LoadChunk(accumulator,   chunk)            ; the running accumulator chunk
    mul   = MulFOp(emb, bcast)                         ; src ln 250
    a_new = AddFOp(mul, a_cur)                         ; src ln 252
            StoreChunk(accumulator, chunk, a_new)      ; inherits AddFOp location
```text

`BroadcastScalarToVector` 以逐 id gain 调用一次（`lowering_util::BroadcastScalarToVector`，`0x13d94460`）；得到的 vector 是每个 chunk 的第二个 `MulFOp` 操作数。这就是完整的逐 combiner weight 应用：把（已经过 combiner 缩放的）gain broadcast 一次，然后逐 chunk fused-multiply-add 进 accumulator。如果存在 `mean`/`sqrtn` divisor，它在到达 HBM 前已经烘入 `g_f32`；此 loop 对 combiner 无感知。

| Op | `lowering_util` symbol | 作用 |
|---|---|---|
| `UnalignedLoadScalarFromHbm` |（逐 id scalar load）| 从 `sorted_gains` 加载原始 32-bit gain word |
| `arith::BitcastOp` | `getF32Type` + `BitcastOp::create` | 把 i32 word 重新解释为 f32 gain |
| `BroadcastScalarToVector` | `0x13d94460` | 把 scalar gain fan 到 lane-width vector（hoisted） |
| `arith::MulFOp` | `MulFOp::create`，`FastMathFlags=none` | `emb_chunk · gain` |
| `arith::AddFOp` | `AddFOp::create`，`FastMathFlags=none` | `+ accumulator_chunk` |
| `LoadChunk` / `StoreChunk` | `0x13d97620` / `0x13d99960` | 读取 embedding/accumulator chunk；写回 |

---

## Inner-Loop 发射

发射结构是三层嵌套。从上到下：逐 sample-tile 的 `scf::ForOp`（在 `Emit` 中）、遍历 sample 的 CSR segment 的逐 id `scf::ForOp`（在 `EmitValencyLoop` 中），以及遍历 feature dimension 的逐 chunk `scf::ForOp`（在 `EmitVectorizedLoop` 中）。每一层的 body 都无条件调用下一层。

### Level 0 — `Emit`：sample-tile loop 和 bounds guard

`Emit`（`0x1332bda0`）构建 sample tile 外层循环，以及丢弃最后一个部分 tile 中越界 lane 的 guard：

```text
Emit outer loop (0x1332bda0)
  if !Target::SupportsSparseCore(): FATAL("SparseCore is not supported by this target")  ; target.h:1709
  lane_cfg = Target[0x948]->[0x94]                       ; per-target max-row / lane config
  n_samples = operand(0).dim                             ; read before first loc tag
  (operand(1).dim mod lane_cfg) RetCheck                 ; src ln 67 "kFeatureWidth % kChunkSize == 0" (remainder != 0 → fail)
  lo   = ConstantIndexOp(0)                              ; src ln 69
  hi   = ConstantIndexOp(n_samples)                      ; src ln 71
  step = ConstantIndexOp(this[+0x60] = ctor-arg n)       ; src ln 73 — tile count (ForOp STEP)
  scf::ForOp(lo, hi, step):                              ; src ln 78 — per-sample-tile loop
    tid     = TileIdOp()                                 ; src ln 86 — SC physical tile id
    s_index = AddIOp(loop_iv, tid-derived)               ; src ln 88 — global sample index
    in_range= CmpIOp(ult, s_index, n_samples)            ; src ln 90 — predicate 6 (ult)
    scf::IfOp(in_range, then = $_0, else = ∅):           ; src ln 95 — bounds guard
      then: EmitSampleCombiner(...) ; YieldOp            ; src ln 98 (lambda $_0 @0x1332ea80)
  SfenceOp("all", 3) ; InsertTileBarrier                 ; src ln 103 — inter-tile synchronization

scf::IfOp 的 then region 运行 EmitSampleCombiner；else region 为空（IfOp 创建时使用 null else builder）。因此这个 IfOp 是tile bounds guard，不是双路 dispatch：它对 in-range sample 执行 combiner，对最后一个部分 tile 的 padding lane 不做任何事。loop 之后，SfenceOp::create(…, "all", 3) 和 lowering_util::InsertTileBarrier 提供 tile 间同步。

NOTE — bounds-guard IfOp 不是 minibatch dispatch。 这个 IfOp 不在 minibatch 和 non-minibatch combiner 之间选择；它只按 s_index < n_samples gate EmitSampleCombiner。non-minibatch combine 路径是独立的 GatherMulScatterSparseDenseMatmulOpDecomposer HLO decomposition，不是这里的 else-branch。

CmpIOp predicate 值（6 = unsigned-less-than）作为内联 register 参数发射，在反编译 C 中不可见为字面量；它来自反汇编读取。predicate 身份为 HIGH；CmpIOp op 及其位置为 CONFIRMED。

Level 1 — `EmitSampleCombiner`：逐 sample accumulator scope

EmitSampleCombiner（0x1332c640）分配一个 scoped f32 SPMEM accumulator，把它清零，运行 valency loop，然后把结果 drain 到 activation output：

text

EmitSampleCombiner (0x1332c640)
  AllocaScopeOp::create                                 ; scope the accumulator region
  F32 = getF32Type
  acc = AllocateScopedMemory(F32, memspace=2)           ; the SPMEM accumulator tile
  InitializeTileSpmemBuffer(acc, 0)                      ; → ZeroMemOp — zero the accumulator
  EmitValencyLoop(builder, acc, ...)                    ; accumulate over the sample's ids
  InitiateSynchronousStreamOperation(...)               ; drain/scatter acc → HBM activation row
  AllocaScopeReturnOp::create                           ; yield from the scope
```text

accumulator 由 `memref::AllocaScopeOp` 限定作用域，因此其生命周期精确为一个 sample；`AllocateScopedMemory` 请求内存空间 `2`（SPMEM tile space）中的 `f32`；`InitializeTileSpmemBuffer` 将其清零（lower 为 `ZeroMemOp` + chunk-iterator loop）。valency loop 填充 accumulator 之后，最终的 `InitiateSynchronousStreamOperation` 把它 scatter 回 HBM 中的 dense activation output，`AllocaScopeReturnOp` 关闭作用域。

### Level 2 — `EmitValencyLoop`：逐 id loop

`EmitValencyLoop`（`0x1332cee0`，四个函数中最大，约 4.8 KB）加载 sample 的 CSR valency，然后遍历该 CSR segment 中的 id，gather 并 FMA 每个 id：

```text
EmitValencyLoop (0x1332cee0)
  lo      = ConstantIndexOp(0)                           ; src ln 151 — loop lower bound
  step    = ConstantIndexOp(1)                           ; src ln 152 — loop step
  valency = UnalignedLoadScalarFromHbm(csr_row_ptr)     ; segment length (no .cc line tag)
  v_idx   = IndexCastOp(valency)                         ; src ln 158 — to index type
  base    = ConstantIndexOp + MulIOp; ConstantIndexOp + MulIOp  ; src ln 160/161 + 164/165 — per-id base offset
  inner   = AllocaScopeOp + getF32Type + AllocateScopedMemory   ; lowering_util_alloc.h:71
  scf::ForOp(0, v_idx, 1) iter_arg = acc:                ; loop over ids in the CSR segment
    id      = AddIOp(loop_iv, ...)                       ; advance the id index
    g_i32   = UnalignedLoadScalarFromHbm(gains_base, id) ; per-id gain word (no .cc line tag)
    g_f32   = BitcastOp(f32, g_i32)                      ; reinterpret as float
    tok_off = ConstantIndexOp; MulIOp; AddIOp            ; per-id embedding token offset
    InitiateSynchronousStreamOperation(tok_off)          ; GATHER this id's embedding row
    EmitVectorizedLoop(acc, g_f32, gathered_row)         ; FMA: acc += row · gain
  AllocaScopeReturnOp::create

valency loop bound 来自 CSR row pointer 的 UnalignedLoadScalarFromHbm（sample 的 segment length），经过 index-cast 后用作 scf::ForOp upper bound。accumulator 作为 loop iter_arg 穿过循环。body 内部，每个 id 加载自己的 gain（bit-cast f32）、计算 embedding token offset、发出同步 gather（由该 offset keyed；InitiateSynchronousStreamOperation 是 stream gather/scatter MRB/FIFO placement 所消费的 LinearStreamStartOp 的生产者），并调用 EmitVectorizedLoop 把 gather 到的 row FMA 进 accumulator。

NOTE — 此 loop 消费的 multiplicity / valency 是 CSR segment length。 每个 sample 的 id 数量（valency）是 loop bound 取自的 CSR row-pointer delta；duplicate-id folding（count × gradient multiplicity）是由 Dedup Multiplicity 负责的预处理步骤，不属于此 emitter。这里 emitter 只是迭代 valency 个 id，每个 id 都带有自己的预计算 gain。

注意事项

发射一个 FMA loop；信任 gain。 不要 special-case sum/mean/sqrtn。reduction 总是 acc += emb · gain；divisor 已在上游折叠进 gain。combiner-enum 分支是建模错误，二进制中没有这种分支。
gain 是 bit-reinterpret，不是数值转换。 UnalignedLoadScalarFromHbm 返回整数；arith::BitcastOp（不是 SIToFPOp）恢复 f32。HBM word 是 float 的 bit pattern。在这里做 integer-to-float convert 会破坏每个 gain。
accumulator 位于 SPMEM，按 sample 设定作用域，并且必须清零。 AllocateScopedMemory(f32, memspace=2) + InitializeTileSpmemBuffer（ZeroMemOp）。跳过 zero-init 会把前一个 sample 的 partial sum 泄漏到下一个 sample。
gain broadcast 一次，而不是每个 chunk 一次。 BroadcastScalarToVector 被 hoist 出 chunk loop；同一个 vector 乘以每个 chunk。逐 chunk 重新 broadcast 是 emitter 避免的浪费。
外层 IfOp 是 bounds guard。 then = combiner，else = empty。它丢弃最后一个部分 sample tile 的 padding lane；它不是 minibatch dispatch，也不是 combiner selector。
未完全映射项（HIGH / LOW / INFERRED）。 CmpIOp predicate 值 6=ult（HIGH — 内联 register arg，来自反汇编读取，不是反编译 C 字面量）。combiner divisor folding：把 mean→1/n / sqrtn→1/sqrt(n) 折叠进 sorted_gains（HIGH — emitter 原样应用 gain 为 CONFIRMED；计算 divisor 的 front-end op 不在 libtpu.so 中，为 INFERRED）。使 gather 成为 indirect id-keyed DMA 而非 linear DMA 的 StreamOptions discriminant bits（LOW — feed 给 stream 的逐 id token offset 为 CONFIRMED；struct 未 bit-decode，见 stream gather/scatter）。EmitSampleCombiner 外层 accumulator 和 EmitValencyLoop 内层 AllocateScopedMemory 是同一个重新设定作用域的 SPMEM buffer 还是两个 allocation（LOW — 二者都经过 memref::AllocaScopeOp；未追踪 aliasing）。GatherMulScatterSparseDenseMatmulOpDecomposer body（non-minibatch combine）（LOW — pass + op name 已锚定；其构建的 HLO 未解码）。

Name	Relationship
`SparseDenseMatmulDotCombinerEmitter::Emit`（`0x1332bda0`）	sample-tile loop + bounds-guard `IfOp`（level 0）
`::EmitSampleCombiner`（`0x1332c640`）	逐 sample SPMEM accumulator scope + zero-init + drain（level 1）
`::EmitValencyLoop`（`0x1332cee0`）	逐 id CSR loop：gain load、gather、FMA call（level 2）
`::EmitVectorizedLoop`（`0x1332e1c0`）	逐 chunk FMA core：`BroadcastScalarToVector` + `MulFOp` + `AddFOp` + `StoreChunk`
`LoweringEmitter::EmitSparseDenseMatmulDotCombiner`（`0x131a7ca0`）	构建 2 个 `FoldAllDimensions` operands 并运行 `Emit` 的 dispatcher
`lowering_util::BroadcastScalarToVector`（`0x13d94460`）	把逐 id gain fan 到 lane-width vector
`lowering_util::{LoadChunk,StoreChunk}`（`0x13d97620` / `0x13d99960`）	FMA 融合跨越的 SPMEM chunk read/write
`lowering_util::InitializeTileSpmemBuffer`（`0x13d93440`）	清零 accumulator（`ZeroMemOp` + chunk-iterator loop）
`lowering_util::InitiateSynchronousStreamOperation`（`0x13d896a0`）	逐 id gather + 逐 sample drain（`LinearStreamStartOp` + `StreamWait` + `SetSyncFlag`）
`GatherMulScatterSparseDenseMatmulOpDecomposer`（`0x1306d740`）	non-minibatch combine counterpart
`XlaSparseDenseMatmulCustomCombinerOnTcWithCsrInputOp`（`0xe653640`）	user-reduce combiner family（替换固定 FMA）

交叉引用

EmitValencyLoop — 孤立的逐 id loop；本页 level 2 构建的 CSR-segment iteration。
Dedup Multiplicity — 产生此 combiner 所迭代 valency 的 duplicate-count / uniquify 预处理。
Embedding Minibatching — 此 lowering 之上的 HLO minibatching decomposition；sample-tile count 的来源。
Stream Gather / Scatter — 此 emitter 对每个 id 发出的同步 indirect gather，以及 drain accumulator 的 scatter。
VectorExtended (VEX) — embedding 可以 lower 到的另一种跨 lane scan/reduce datapath；DotCombiner 是 gather-FMA 形式，不是这个路径。
Scan Datapath — DotCombiner 不使用的 SegmentedScanOp reduce 路径。
SparseCore Overview — 三个 SC engine class，以及 embedding gather-reduce-scatter datapath 所处位置。
Binary: extracted/libtpu-0.0.40-cp314-cp314-manylinux_2_31_x86_64/libtpu/libtpu.so（build-id 89edbbe81c5b328a958fe628a9f2207d）
Index entry: Part IX — SparseCore & BarnaCore / SparseCore datapath (embeddings) — back to index

SampleCombiner 发射器 ​

摘要 ​

Combiner Reduction 名册 ​

一个 reduction，三个名称 ​

Broadcast 并应用 gain ​

Level 1 — EmitSampleCombiner：逐 sample accumulator scope ​

注意事项 ​

相关组件 ​

交叉引用 ​

SampleCombiner 发射器

摘要

Combiner Reduction 名册

一个 reduction，三个名称

Broadcast 并应用 gain

Level 1 — `EmitSampleCombiner`：逐 sample accumulator scope

注意事项

相关组件

交叉引用