EmitValencyLoop

本页中的每个地址、op-create 顺序、scf::ForOp bound、stack-slot identity 和 ::create 签名都逐字节精确地读取自 libtpu-0.0.40-cp314 wheel 中的 libtpu.so（build-id 89edbbe81c5b328a958fe628a9f2207d；构建 libtpu_lts_20260413_b_RC00）。.text VA 等于文件偏移 0xe63c000；.rodata 位于 0x84a0000；二者都是 identity-mapped。该二进制未 strip，nm -C 可解析每个方法。EmitValencyLoop 是单一的、generation-agnostic 的 xla::tpu::sparse_core emitter；它恰好只有一个实例，在运行时由 xla::jellyfish::Target& 参数化，而不是每代（vfc/glc/gfc）命名空间 clone。本页地址均来自此构建。

摘要

SparseDenseMatmulDotCombinerEmitter::EmitValencyLoop (@0x1332cee0, 0x12e0 B) 是 per-sample 标量 combiner lowering 的最内层 embedding 循环。对于一个 sample（一个 CSR 行），它在 HBM-load 时加载该 sample 的 valency，即该 sample 的 CSR segment 中 id 的运行时数量；构建一个 trip count 就是该 valency 的 scf::ForOp；并在循环内部为每个 id gather 一行 embedding（同步间接 stream），再通过 EmitVectorizedLoop 将 emb·gain FMA 到 SPMEM accumulator。它是 embedding sum-lookup 的标量 per-id 形式。

本页记录三件事，均已对照该函数的反编译函数体确认：(1) 它生成的三级循环结构，即 per-sample setup、per-id scf::ForOp 和 per-feature-chunk vectorized FMA，以及它用 memref::AllocaScopeOp region 作用域化的内部 SPMEM accumulator；(2) valency 机制，即驱动 for upper bound 的运行时 UnalignedLoadScalarFromHbm → IndexCastOp，以及有意不存在任何运行时 valency division（mean / sqrtn 在前端折叠进 per-id gain）；(3) 此循环旁边的 sort / uniquify / dedup 接线，即 HLO SparseDenseMatmulOpDecomposer dedup datapath（SortAndPartition → ReduceDuplicates CSR→ELL → 对 unique ids 做 SparseGather → segmented reduce）以及这些 lowering 到的 SC-dialect SortOp / UniqueOp / UniqueWithLaneIdsOp / SegmentedScanOp / DuplicateCount op 集合。DotCombiner per-id 标量循环和 Sort/Unique dedup 路径是同一个 embedding sum-lookup 的两种替代 lowering；本页锚定 valency loop 如何接线，以及 dedup 操作数如何穿行。

对重新实现而言，契约是：

Valency 是运行时标量，不是编译期常量。 EmitValencyLoop 针对 CSR/id 操作数发出 UnalignedLoadScalarFromHbm，将结果 IndexCastOp 到 index，并把它用作 scf::ForOp upper bound。生成的循环是 for (i = 0; i < valency; ++i)。
循环中任何位置都没有运行时除法。 穷尽地说：没有 arith::DivFOp、math::RsqrtOp、reciprocal、pow，也没有任何 divsd/vdiv/vsqrt 指令。mean (1/n) 和 sqrtn (1/√n) combiner 除数由 TF/JAX 前端预先折叠进 per-id gain。valency 只驱动整数 trip count。
每个 id：先 gather，再 FMA。 函数体加载 per-id gain（UnalignedLoadScalarFromHbm → arith::BitcastOp i32→f32），计算 per-id token offset（feature_width × 4 byte-scaled，MulIOp + AddIOp），发起同步间接 gather（InitiateSynchronousStreamOperation），并调用 EmitVectorizedLoop(gather_buf, gain, outer_accumulator)。
Accumulator 位于 AllocaScopeOp region 中。 memref::AllocaScopeOp 打开内部 region；AllocateScopedMemory 在其中分配 f32 SPMEM accumulator；memref::AllocaScopeReturnOp 在循环结束时 yield 出结果。
Dedup 路径是另一个 dedup-optimised lowering。 Sort/Unique/ReduceDuplicates/SparseGather datapath 属于 SparseDenseMatmulOpDecomposer（HLO custom-call decomposition），不属于此 emitter。此 emitter 是 per-id 标量 gather-FMA 形式。


函数	`xla::tpu::sparse_core::SparseDenseMatmulDotCombinerEmitter::EmitValencyLoop(OpBuilder, Value v1, Value v2)` `@0x1332cee0` (`0x12e0` B)
调用方	`…::EmitSampleCombiner` `@0x1332c640`（call site `@0x1332ca82`）
`v1`	CSR/embedding-id 操作数 `Value`，即 `UnalignedLoadScalarFromHbm` 源（valency + gain）
`v2`	SPMEM outer accumulator `Value`（刚被 `InitializeTileSpmemBuffer` 清零的 tile）
循环原语	`scf::ForOp::create` `@0x17866d60` — `(lower=ConstIdx(0), upper=IndexCast(valency), step=ConstIdx(1), iter_args={acc})`
Valency load	`lowering_util::UnalignedLoadScalarFromHbm` `@0x13da4580` (ComputeType=1, operand=`v1`) → `IndexCastOp` `@0x1cb0ce80`
Per-id gather	`lowering_util::InitiateSynchronousStreamOperation` `@0x13d896a0` (edx=1, r8d=1, `StreamOptions` by value)
Inner FMA	`EmitVectorizedLoop` `@0x1332e1c0` `(v1=gather_buf, v2=gain, v3=outer_acc, v4=unset)`
Accumulator scope	`memref::AllocaScopeOp::create` `@0x18304960` + `AllocaScopeReturnOp::create` `@0x18304e40`
Division ops	无 — 函数体中没有 `DivFOp`/`RsqrtOp`/reciprocal/`sqrt`/`divsd`/`vdiv`
源文件	`platforms/xla/sparse_core/sparse_dense_matmul_dot_combiner_emitter.cc` (`@0x87610c0`)

循环结构

目的

EmitValencyLoop 构建 per-sample → per-id → per-feature-chunk 嵌套，把一个由 (id, gain) pair 组成的 CSR 行转换为一个累积的 embedding-sum 向量。它自身不会在编译期迭代；它会发出 MLIR op，即 scf::ForOp、memref::AllocaScopeOp、arith op 和两个嵌套 emitter 调用，随后 SC 后端把它们降低为 sequencer 程序。下面的结构是此次 emission 的静态 op-create 顺序。

签名与 ABI

EmitValencyLoop(OpBuilder, Value v1, Value v2)，this 位于 rdi，OpBuilder 以值传递于 [rbp+0x10]。调用方 EmitSampleCombiner（call site @0x1332ca82）传入：

实体	Reg / slot（入口）	身份
`this`	`rdi` → `rbx`, `[rbp-0x30]`	`SparseDenseMatmulDotCombinerEmitter*`（`+0x8` = builder/context state，`+0x10` = HLO，`+0x38` = `Target&`）
`v1`	`rsi` → `r13`	CSR/embedding-id 操作数 `Value`（`EmitSampleCombiner` `[rbp-0xc8]`），即 `UnalignedLoad` 源
`v2`	`rdx` → `r14`, `[rbp-0x130]`	SPMEM outer accumulator `Value`（已清零 tile；作为 `v3` 穿入 `EmitVectorizedLoop`）
`OpBuilder` (by value)	`[rbp+0x10]` (`r12`)	插入 builder；ymm-copy 到 stream / loop 调用帧

NOTE — v2 已由调用方预先清零。 EmitSampleCombiner 在把 tile 作为 v2 传入前，会对其调用 InitializeTileSpmemBuffer（@0x13d93440，call site @0x1332ca53）；EmitValencyLoop 会累加到其中，而不会清零它。

Op-Create 顺序（逐字节精确）

该函数恰好发出以下 op 序列。每行都是列出 VA 处的一次 …::create 调用；src ln 是在每个 op 之前的 DebugInfoTracker::GetCurrentLocation 调用中烘入的 .cc 行号。

text

#   op / call                                   @VA          src ln  role
--  ------------------------------------------   ----------   ------  --------------------------------------------
0   HloInstruction::operand(1)  ×2               1332cf06 /   151     read CSR/id operand-1 Shape twice; dims[0],
                                                 1332cf4b             dims[1]; layout CHECK (fail @1332e13d ln843)
1   arith::ConstantIndexOp(0)                    1332d038     151     ForOp LOWER  → [rbp-0x118]
2   arith::ConstantIndexOp(1)                    1332d0f4     152     ForOp STEP   → [rbp-0xf0]
3   UnalignedLoadScalarFromHbm (CSR valency)     1332d185     160     load per-sample valency; operand = v1
4   arith::IndexCastOp (valency → index)         1332d247     —       → ForOp UPPER [rbp-0xe8]
5   ConstantIndexOp + arith::MulIOp #1 (hi dim)  1332d3b2 /   164/165 CSR row addr from operand1.dim0 (hi half)
                                                 1332d3c7
6   ConstantIndexOp + arith::MulIOp #2 (lo dim)  1332d515 /   164/165 CSR row base → [rbp-0x148]
                                                 1332d52a
7   TypeRange + memref::AllocaScopeOp::create    1332d621     71      open inner-accumulator region (lowering_util_alloc.h)
8   getF32Type + AllocateScopedMemory            1332d792 /   —       f32 SPMEM inner accumulator → r12
                                                 1332d7eb
9   scf::ForOp::create  (VALENCY LOOP)           1332d8d4     —       (Const0, IndexCast(valency), Const1, {acc})
--- loop body (per id) ---------------------------------------------------------------------------------------
10  arith::AddIOp #1 (id index)                  1332da10     —       loop_iv + CSR base from MulIOp [rbp-0x148]
11  UnalignedLoadScalarFromHbm (per-id gain)     1332da93     195     load per-id sorted_gains scalar; operand = v1
12  getF32Type + arith::BitcastOp                1332db43 /   198     i32 bits → F32 gain → [rbp-0xe8]
                                                 1332db54
13  (feat «×4») + ConstantIndexOp + MulIOp #3    1332dcb9 /   —       per-id token offset
                                                 1332dd82
14  arith::AddIOp #2 (token offset + base)       1332dd9b     —       final embedding-row offset
15  InitiateSynchronousStreamOperation           1332defb     —       indirect gather (edx=1, r8d=1, StreamOptions)
16  EmitVectorizedLoop call                      1332df2e     —       FMA acc += emb·gain
--- end loop body --------------------------------------------------------------------------------------------
17  memref::AllocaScopeReturnOp::create          1332e08d     —       yield accumulator out of inner scope
18  ret                                          1332e13c     —       StatusOr<>; CHECK-fail path @1332e13d ln843
```text

> **GOTCHA — operand(1) Shape 以 3-vs-5 storage discriminant 读取。** 两次 `operand(1)` 读取（`@0x1332cf06`, `@0x1332cf4b`）都会从 `[hlo+0x58]` 加载该操作数的 `Shape*`，然后按 `[shape+0x138]`（offset 312）处的字节分支：`== 3` → dims inline；`== 5` → dims on heap，先解引用 `[shape]`；其他任何值 → `LogMessageFatal("buffer != nullptr", shape.h:843)`。重新实现者在读取 CSR/id 操作数的 dims 时必须复现这个 inline/heap small-buffer discriminant，否则 layout 读取会错误。`dims[0]`（`[…+0x8]`）→ `r15` 是 per-id stride；`dims[1]`（`[…+0x10]`）→ `[rbp-0xe0]` 是 feature width。

### 三级嵌套

```text
region                         ops emitted in it                                  role
-----------------------------  ------------------------------------------------   --------------------------
EmitValencyLoop body           operand(1) shape; Const0/Const1; CSR-valency       set up loop bounds, the CSR row
  (outer, per sample)          load + IndexCast; CSR-addr MulIOp×2; AllocaScopeOp; base, and the inner accumulator
                               AllocateScopedMemory (inner acc); scf::ForOp
scf::ForOp body                AddIOp(id index); gain load + Bitcast(i32→f32);     per-id gather + FMA
  (per id; iter_arg = acc)     token-offset MulIOp/AddIOp; Initiate-              (one embedding row per id)
                               SynchronousStreamOperation (gather);
                               EmitVectorizedLoop (FMA into acc)
EmitVectorizedLoop body        BroadcastScalarToVector(gain); LoadChunk×2;         per-feature-chunk FMA
  (per feature chunk)          MulFOp(emb·gain); AddFOp(+acc); StoreChunk(→acc)

外层 region 是函数体本身。中层 region 是 scf::ForOp，其唯一 iter_arg 是 SPMEM accumulator。最内层 region 由 EmitVectorizedLoop 生成，它把标量 gain broadcast 成向量，把 gather 到的 embedding row 乘以它，并逐 chunk 加到 accumulator 中。

Valency 机制

目的

Valency 是参与一个 sample lookup 的 embedding id 数量；它随 sample 变化，只有运行时才知道。本单元记录 EmitValencyLoop 如何读取它、它如何成为 loop bound，以及一个关键的负面结果：绝不会对它应用 division。

运行时 Valency Load → `scf::ForOp` Bound

per-sample valency 在运行时从 CSR row buffer 加载，不会在编译期烘入：

// 由 EmitValencyLoop 发出（op-create #3, #4, #9）
ScalarFromHbm = UnalignedLoadScalarFromHbm(/*Target&*/, /*ComputeType=*/1,
                                           /*OpBuilder*/, /*SmallVector<Value,6>=*/{v1});  // @0x1332d185
valency_idx   = arith::IndexCastOp(loc, IndexType, ScalarFromHbm);                          // @0x1332d247
loop = scf::ForOp::create(builder,
                          /*lower=*/ConstantIndexOp(0),        // [rbp-0x118]
                          /*upper=*/valency_idx,               // [rbp-0xe8]  ← 运行时 valency
                          /*step =*/ConstantIndexOp(1),        // [rbp-0xf0]
                          /*iter_args=*/ValueRange{accumulator});   // @0x1332d8d4
```text

`SmallVector<Value,6>` 参数以内联方式构造，header word 为 `0x600000001`（capacity 6，size 1），持有 `{v1}`，两个 load site 均已确认。`ComputeType` 两处都以 `esi = 1` 传递。因此生成的循环是 `for (i = 0; i < valency; ++i)`，SPMEM accumulator 作为 loop-carried value 穿行。

> **NOTE — valency 是从 HBM 读取的 CSR segment length。** `v1`（CSR/embedding-id 操作数）是 load 源；加载的标量是 per-sample row pointer / id-count，经 `IndexCast` 后驱动 trip count。它不是常量，不同 sample 会迭代不同次数。

### Per-Id Gain Load（不是 Division）

在循环内部，per-id gain 作为原始 i32 位加载，并 **bitcast** 到 f32；它从不通过 division 计算：

```c
// op-create #11, #12 — 位于 scf::ForOp body 内
gain_bits = UnalignedLoadScalarFromHbm(/*ComputeType=*/1, /*SmallVector=*/{v1});  // @0x1332da93 (src ln 195)
gain_f32  = arith::BitcastOp(loc, getF32Type(), gain_bits);                       // @0x1332db54 (src ln 198)

gain_f32 作为 v2 传给 EmitVectorizedLoop，在那里它变成 BroadcastScalarToVector(gain) → MulFOp(emb·gain) → AddFOp(+acc)。因此 combiner 是 weighted-sum FMA；除数存在于所提供的 gain 内部。

QUIRK — 两次 UnalignedLoadScalarFromHbm 调用共享操作数 v1 和 ComputeType=1；offset 区分 valency 与 gain。 CSR-valency load（@0x1332d185）和 per-id gain load（@0x1332da93）传递相同的 SmallVector<Value,6>={v1} 与 ComputeType=1；只有计算出的 token offset 不同。这与 v1 是一个同时承载 CSR row pointer 和 sorted gains 的单一 packed/strided HBM buffer 相一致，并由不同 offset 索引；不过单 buffer multiplexing 尚未被位级证明，gain 也可能是通过同一 Value 的不同 base 到达的不同 buffer。

没有运行时 Valency Division — 穷尽确认

EmitValencyLoop 中完整的 …::create 调用集合是：ConstantIndexOp ×5、MulIOp ×3、AddIOp ×2、IndexCastOp ×1、BitcastOp ×1（i32→f32 gain reinterpret）、AllocaScopeOp / AllocaScopeReturnOp、valency scf::ForOp、UnalignedLoadScalarFromHbm ×2、InitiateSynchronousStreamOperation ×1，以及 EmitVectorizedLoop ×1。函数体中没有 arith::DivFOp、math::RsqrtOp、reciprocal、PowOp，也没有任何 sqrt/divsd/vdiv/vsqrt/vrsqrt x86 指令。

QUIRK — mean / sqrtn divisor 是前端 gain-scale，不是运行时 op。 对于 mean combiner，除数是 1/n；对于 sqrtn，它是 1/√n。两者都不在这里计算：前端（TF/JAX）在 gain 到达 HBM 前把除数折叠进 per-id gain，所以 EmitValencyLoop 原样应用 gain（load → bitcast → 在 FMA 中相乘）。valency 只用于确定整数 scf::ForOp 大小。重新实现者必须复现 gain-side folding；期望在这里找到运行时 divide 是找不到的。

Per-Id Gather 与 FMA 接线

目的

每次循环迭代恰好从 HBM gather 一行 embedding 到 SPMEM，并将其 FMA 到 accumulator 中。本单元钉住作为 gather key 的 token-offset arithmetic，以及 gather 与 FMA 的调用参数 identity。

Token-Offset Arithmetic

per-id embedding-row offset 由 feature width（operand1 dims[1]）和 loop index 构建：

// op-create #10, #13, #14 — gather index
id_index  = arith::AddIOp(loop_iv, csr_base);          // @0x1332da10  (csr_base from MulIOp #2 [rbp-0x148])
feat_x4   = feature_width << 2;                         // @0x1332dcb9  (×4 = i32 bytes per word)
tok_mul   = arith::MulIOp(ConstantIndexOp(feat_x4), …); // @0x1332dd82  (MulIOp #3)
token_off = arith::AddIOp(tok_mul, base);              // @0x1332dd9b  → the embedding-row offset
```text

`feature_width << 2`（反编译中的 `v223 = 4 * v223`，`@0x1332dcb9`）把 feature count 转换为 4 字节（i32）word 的 byte stride。

### 同步间接 Gather

```c
// op-create #15
InitiateSynchronousStreamOperation(/*Target&*/, /*Array=*/token_off_operands,
                                   /*edx=*/1, /*r8d=*/1, /*StreamOptions by value*/);  // @0x1332defb → @0x13d896a0

这会发起该 id 的 embedding row 的间接 HBM→SPMEM gather，以 token_off 为 key。两个前导标量参数都是 1（edx=1, r8d=1）；它们是 StreamOptions 判别位，用于选择indirect id-keyed load，而不是线性 DMA。per-id token offset 是 gather index；两个 StreamOptions bool 的精确含义未做位级解码。这个 stream 的 producer/consumer 侧是 Stream Gather/Scatter 页上的 LinearStreamStartOp / IndirectStreamStartOp 家族。

`EmitVectorizedLoop` 调用接线

text

arg     reg / slot           identity                          confirmed via
------  ------------------   -------------------------------   ------------------------------------------
this    rdi = rbx            the emitter                        member layout
builder [rbp-0x110]          inner-region OpBuilder (in loop)   ymm copy to [rsp]
v1      rdx = [rbp-0xf0]     the gathered embedding / inner buf set from [rbp-0xf0] at 1332dde6
v2      rcx = [rbp-0xe8]-0x10 the GAIN (Bitcast result)         BroadcastScalarToVector(v2) = gain
v3      r8  = [rbp-0x130]    the OUTER accumulator (= v2 saved) LoadChunk / StoreChunk(v3) = acc
v4      (r9 unset)           unused / null Value                4th Value left unset
```text

调用（`@0x1332df2e`）目标为 `EmitVectorizedLoop @0x1332e1c0`。其 mangled name 携带四个 `Value` 参数（`…N4mlir9OpBuilderENS3_5ValueES5_S5_S5_`），确认是 4-`Value` 签名；第 4 个（`r9`）未加载，因此 `v4` 作为默认/null `Value` 传入。在两个 fallible 调用之后，`StatusOr` ok-check（`cmp rax, 1`）控制 success path；failure path 在 `.cc` lines 208 / 210 构建 `absl::Status`。

### Drain

`memref::AllocaScopeReturnOp::create`（`@0x1332e08d` → `@0x18304e40`）将 accumulator 结果 yield 出 `AllocaScopeOp` region；`mlir::ValueRange::ValueRange` 打包该 op 的结果，函数在成功时以 `*(_QWORD*)ptr = 1` 返回 `StatusOr<…>`。

---

## Sort / Uniquify / Dedup 接线

### 目的

`EmitValencyLoop` 是 embedding sum-lookup 的 per-id **标量**形式。还存在另一种 *dedup-optimised* lowering：HLO `SparseDenseMatmulOpDecomposer` decomposition 会排序 id、折叠重复项（CSR→ELL）、在 unique window 上 gather，并做 segment-reduce。本单元记录这条接线，即此循环旁边的层，以及它 lowering 到的 SC-dialect op 操作数/结果穿线。DotCombiner emitter 本身不运行 sort/unique DAG；dedup 路径属于 decomposer。

### HLO Dedup Datapath (`SparseDenseMatmulOpDecomposer`)

```text
stage              function                                       @VA          role
-----------------  --------------------------------------------   -----------  --------------------------
sort + partition   SortByPartitionIdsOrFallthrough                1366ea60     sort (sample,token) ids
                   SortAndPartitionSparseInput                    1366f9c0       "
                   PartitionCsr / PartitionCsrByLogicalReplicas   13670100 /   per-physical-core CSR window
                                                                  1366fae0
reverse index      CalculateReverseIndex                          13670f00     un-permute index
DEDUP + reduce     ReduceDuplicates                               136722e0     CSR→ELL duplicate collapse
gather (unique)    DistributedGatherUniqueVectors                 13674080     gather over unique ids
                   GatherAndQuantizeEmbeddingVectors              13674680       "  + dequant
                   GatherEmbeddingVectors                         13674d60       "
reduce             ReduceVectors / FixedWindowReduceVectors       13676680 /   per-sample reduce (× dup count)
                                                                  13675d20
activations        ComputeActivations                             13675840     final embedding-sum output

该链条镜像更广泛的 embedding-minibatching decomposition 中的 gather→sort→uniquify→scan→scatter DAG；上面的地址是 forward-embedding decomposition 入口点。

`ReduceDuplicates` — CSR→ELL Dedup Op

ReduceDuplicates (@0x136722e0) 是 duplicate-collapse + segment-reduce。其 shape contract 由 SparseCoreShapeVerifier::HandleReduceDuplicates (@0x136538a0) 强制执行：reduce->operands().size() == 2（"ReduceDuplicates expects two inputs"，即 sorted-CSR input + init value）、reduce->called_computations().size() == 1（"expects a reduce function"，即 combiner，例如 add），以及 "Csr input and Ell output must have identical number of rows"（CSR→ELL collapse）。其函数体发出：

text

#   op / custom-call                              @VA          notes
--  -------------------------------------------   ----------   -----------------------------------------
0   CreateGetTupleElement ×2 + CreateTuple(2)      136722.. /   unpack {csr_pointers, values}, re-tuple
                                                   13672573
1   CustomCall "SparseMapRow"  (op-type 0x4)       1367260b     5-arg form, WITH the reduce HloComputation
2   CreateGetTupleElement ×2                        13672897 /   extract SparseMapRow results
                                                   13672930
3   reduce computation: CreateParameter ×3          13672a70     lhs/rhs value params + init; nested SparseMapRow
4   CustomCall "DynamicBoundedSlice" (op-type 0x11) 13672e56     bound the unique window to max_unique_ids
5   CreateTuple + CustomCall "SparseMapRow" (0x4)   13672fe8     final map-row
6   CreateGetTupleElement + AddInstruction          13673588     package
```text

op-type 名称从 `SparseCoreOperationTypeToString` jump table（`@0x14b7f480`，表 `@0xaf36d70`）逐字节解码：entry `0x4` = `"SparseMapRow"`，entry `0x11` = `"DynamicBoundedSlice"`。duplicate multiplicity 由 `SparseMapRow` 的 per-ELL-row reduce 实现，它会折叠已排序的重复 token-id；`DynamicBoundedSlice` 将 deduped window 限制到配置 `max_unique_ids_per_partition`。

> **NOTE — dedup 直接缩小 gather。** `DistributedGatherUniqueVectors`（`@0x13674080`）lowering 到 HLO `SparseGather` op，其契约（`SparseCoreShapeVerifier::HandleSparseGather @0x1365b880`）是："SparseGather expects 2 operands"（embedding table + unique-id gather indices）、"require a CSR Config"，以及 "Gathered Output dimension-0 size must match input CSR's max-non-zeros"。输出 dim-0 是 *unique-id* window，所以对 id 去重会直接减少 embedding-table HBM gather 的数量。这就是 redundant-gather elimination，与 `EmitValencyLoop` 的 one-gather-per-id 标量形式相对。

### SC-Dialect Op 签名（由 `::create` 逐字节确认）

上述 HLO op lowering 到 SC dialect。`::create` 签名从（未 strip 的）symbol mangling 和 producer call site 确认：

| Op | `::create` @VA | 签名 | Producer 调用点 |
|---|---|---|---|
| `SortOp` | `0x14604600` (Type form) / `0x14604800` (TypeRange) | `(OpBuilder, Location, Type, Type, Type, Value, Value, Value, StringAttr)` — **3 result types, 3 key inputs, sort-dir attr** | `FusionEmitter::SetOutputWindowBoundsForSparseMapRow` (`@0x1389136a`) → `SortOpLowering::matchAndRewrite @0x13597700` → `RadixSortEmitter` |
| `UniqueOp` | `0x14622400` (also `0x14622200`) | `(OpBuilder, Location, Value, Value)` — **2 sorted-key inputs** → unique values + index/marker | `RankAndPermuteComputeFunction` (`@0x134042a9`) — `rdx=key0`, `rcx=key1` |
| `UniqueWithLaneIdsOp` | `0x146231a0` | `(OpBuilder, Location, Value, Value)` — same 2 inputs, **+ per-lane (multiplicity) outputs** (`getNextResultAtOffset` ×2) | `RankAndPermute` (`@0x13404303`) |
| `SegmentedScanOp` | `0x145fd5a0` | `(OpBuilder, Location, Type, Value data, Value seg-id, StringAttr reduction_op)` | `ScanOpLowering @0x135f348d`; `ComputeElementScatter @0x13d2a766` |
| `DuplicateCountOp` / `…WithLaneIds` | (inline build) | lowered via `DuplicateCountUniqueOpLowering<T>` → LLVM struct `{i32 count, f32 …}` → `ReplaceOpWithExtracts` | `matchAndRewrite @0x13599d40` / `@0x1359a7e0` |

`SortOp` 携带一个 4 字符 sort-direction `StringAttr` `"dscd"`（`@0x8720761`，通过 `StringAttr::get` 构建）；1 输入 HLO `SortLexicographic`（"SortLexicographic expects one input"，`HandleSortLexicographic @0x13653540`）经由 `VariableWindowPipelineEmitter` → dialect `SortOp` → `SortOpLowering` → `RadixSortEmitter` lowering。`UniqueOp` 和 `UniqueWithLaneIdsOp` 在 `RadixSortEmitterInternal::RankAndPermuteComputeFunction`（`@0x134039c0`）中被调用于已排序的 `(token, sample)` key column 上。

### Multiplicity → Segmented Reduce

`DuplicateCountOp`、`DuplicateCountWithLaneIdsOp`、`UniqueOp` 和 `UniqueWithLaneIdsOp` 全部共享**同一个** LLVM lowering 模板 `DuplicateCountUniqueOpLowering<T>`，其实例化位于 `0x13599d40` / `0x1359a7e0` / `0x1359b280` / `0x1359bd20`，并且都在 `SortOpLowering` 旁边注册到同一个 `RewritePatternSet::add`（`@0x13572820`）中。该 lowering 构建一个 LLVM struct `{i32, f32, …}`，并通过 `ReplaceOpWithExtracts` 提取结果；其中 `i32` 是**每个 unique id 由多少个已排序 entry 折叠而来**，也就是 per-unique valency weight。

这个 multiplicity 会缩放 segmented reduce。`SegmentedScanOp(resultType, data, seg-id, reduction_op)`（`reduction_op` 通过 `getReductionOp` 读取）在每个 CSR 边界重置；因为 duplicate-collapse 已经为每个 unique entry 预先按 row population 加权，所以 `SegmentedScan` 会在 multiplicity 已烘入的情况下对 unique window 做 reduce。关于 `DuplicateCount`→multiplicity 与 `Uniquify` inverse-permutation 细节，见 [dedup multiplicity](dedup-multiplicity.md)；关于 reduce，见 [segmented add scan](segmented-add-scan.md) / [scan datapath](scan-datapath.md)。

> **NOTE — 两种 lowering，一种语义。** DotCombiner per-id 标量循环（本页）和 Sort/Unique/`ReduceDuplicates`/`SparseGather` dedup datapath 是同一个 embedding sum-lookup 的两种替代 lowering。标量循环更简单（每个 id 一次 gather，gain 原样应用）；dedup 路径消除冗余 gather，但会穿行 sort/unique/segmented-scan op 集合。给定 embedding op 采用哪条 lowering 是 decomposer 的决定，不是 `EmitValencyLoop` 的决定。

---

## Generation 覆盖

`SparseDenseMatmulDotCombinerEmitter::EmitValencyLoop` **没有**按 generation clone。二进制在 generation-neutral 的 `xla::tpu::sparse_core` 命名空间中恰好带有一个定义（`@0x1332cee0`，mangled `_ZN3xla3tpu11sparse_core35SparseDenseMatmulDotCombinerEmitter15EmitValencyLoop…`）；`nm` 找不到第二份副本。Generation 选择是*运行时*参数：外层 `SparseDenseMatmulDotCombinerEmitter` 构造器接受 `xla::jellyfish::Target const&`（`@0x1332bd40`），而它调用的 helper（`UnalignedLoadScalarFromHbm`、`AllocateScopedMemory`、`InitiateSynchronousStreamOperation`）都会在内部按该 `Target&` 分支。因此相同 emitter 和相同 op-create 序列服务于每个带 SparseCore 的 target；没有可用于 diff 的 vfc/glc/gfc 命名空间 sibling。

二进制其他位置存在的 generation-namespaced 家族（`gxc`/`glc`/`gfc`，`vxc`/`vfc`）是此 emitter 之下的 backend/datapath 层，而不是它的副本。SparseCore embedding datapath 在 pre-SparseCore generation（jellyfish=v2，pufferfish=v4）上不存在；`xla::tpu::sparse_core` 命名空间和此 emitter 只出现一次。

---

## 限制与开放项

| 项目 | 备注 |
|---|---|
| `EmitValencyLoop` full op-create order + `scf::ForOp` bounds (Const0, IndexCast(valency), Const1, {acc}) | 从函数体读取 |
| Runtime valency load → `IndexCast` → loop upper bound | 两个 load site + IndexCast 已读取 |
| Absence of any runtime division / sqrt / reciprocal in the loop | 穷尽 op-create + x86 指令扫描 |
| Per-id gain load + i32→f32 `BitcastOp` | accessor + `BitcastOp::create` 已读取 |
| Token-offset arithmetic (`feat << 2`, MulIOp + AddIOp) → gather index | 从函数体读取 |
| `EmitVectorizedLoop` call args (`v1=gather_buf, v2=gain, v3=outer_acc, v4=unset`) | reg/slot + 4-Value mangling |
| `ReduceDuplicates` contract (2 inputs + 1 reduce-fn, CSR→ELL) + `SparseMapRow`(0x4)/`DynamicBoundedSlice`(0x11) | verifier strings + jump table |
| SC-dialect `Sort`/`Unique`/`UniqueWithLaneIds`/`SegmentedScan` `::create` signatures | symbol mangling + producer sites |
| `DuplicateCountUniqueOpLowering<T>` → struct-extract multiplicity | 4 个 template instantiation + pattern-set |
| Whether `v1` is a single packed CSR+gain buffer vs. two buffers via one `Value` | shared operand + ComputeType 已读取；multiplexing 未被位级证明 |
| `StreamOptions` bits (edx=1/r8d=1) selecting indirect vs. linear DMA | per-id offset 是 gather index；bool 未解码 |
| `SortOp` `"dscd"` per-char semantics (direction / stability per key column) | 读取为 4 字符 attr；per-char 含义未交叉解码 |
| Exact `ReduceDuplicates`→`SparseMapRow`→(`SortOp`+`ScanOp`) inlining chain | `SparseMapRow` 携带 reduce fn + 构建 Sort/Scan window；函数体未完全解码 |
| The precise emitter that creates the dialect `DuplicateCountOp` (no standalone `::create`; built inline) | LLVM lowering + `UniqueWithLaneIds` lane outputs 已读取；dialect-create producer 未隔离 |

---

## 交叉引用

- [SampleCombiner Emitter](sample-combiner-emitter.md) — `EmitSampleCombiner` 调用方，以及此循环调用的 `EmitVectorizedLoop` inner FMA。
- [RankAndPermute / RadixSort](rank-and-permute-radixsort.md) — `UniqueOp` / `UniqueWithLaneIdsOp` 消费 sorted keys 的 sort/permute compute function。
- [Dedup Multiplicity](dedup-multiplicity.md) — 为 reduce 加权的 `DuplicateCount` → multiplicity 和 `Uniquify` inverse-permutation。
- [Stream Gather/Scatter](stream-gather-scatter.md) — per-id `InitiateSynchronousStreamOperation` 发起到的 indirect-stream slot。
- [VectorExtended (VEX)](vectorextended-vex.md) — inner per-feature-chunk FMA 运行所在的 vector datapath。
- [Embedding Minibatching Decomposition](embedding-minibatching.md) — 上方 HLO 层，包含完整 gather→sort→uniquify→scan→scatter DAG。
- [Segmented Add Scan](segmented-add-scan.md) / [Scan Datapath](scan-datapath.md) — dedup 输入的 `SegmentedScanOp` per-segment reduce。
- [SparseCore 概览](overview.md) — host-table → HBM → SC embedding datapath，本循环是其中的 per-id 形式。
- **Binary:** `extracted/libtpu-0.0.40-cp314-cp314-manylinux_2_31_x86_64/libtpu/libtpu.so` (build-id `89edbbe81c5b328a958fe628a9f2207d`)
- **索引条目:** Part IX — SparseCore & BarnaCore / SparseCore datapath (embeddings) — [返回索引](../index.md)

EmitValencyLoop ​

摘要 ​

循环结构 ​

目的 ​

签名与 ABI ​

Op-Create 顺序（逐字节精确） ​

Valency 机制 ​

目的 ​

运行时 Valency Load → scf::ForOp Bound ​

没有运行时 Valency Division — 穷尽确认 ​

Per-Id Gather 与 FMA 接线 ​

目的 ​

Token-Offset Arithmetic ​

EmitVectorizedLoop 调用接线 ​

ReduceDuplicates — CSR→ELL Dedup Op ​

EmitValencyLoop

摘要

循环结构

目的

签名与 ABI

Op-Create 顺序（逐字节精确）

Valency 机制

目的

运行时 Valency Load → `scf::ForOp` Bound

没有运行时 Valency Division — 穷尽确认

Per-Id Gather 与 FMA 接线

目的

Token-Offset Arithmetic

`EmitVectorizedLoop` 调用接线

`ReduceDuplicates` — CSR→ELL Dedup Op