Matprep、IAR 和 Latch 子槽

地址适用于 libtpu-0.0.40-cp314 wheel 中的 libtpu.so（BuildID md5 89edbbe81c5b328a958fe628a9f2207d，未 stripped，含完整 C++ 符号）。其他版本会不同。

摘要

在 MXU slot 的脉动阵列能够执行一次乘法时钟步之前，它必须先被喂数。喂给阵列包含三个不同的工作，本页覆盖这些工作；它们都通过矩阵乘法本身使用的同一个 VectorExtended bundle slot 运行：matprep 暂存移动操作数（activations）并准备静态增益矩阵；IAR（Index Address Register）是每个 TensorCore 的索引寄存器文件，用来驱动进入操作数池的 indexed（gather）内存访问；而 latch op 会把准备好的增益矩阵加载到阵列的权重 bank 中。矩阵乘法步骤本身是可见的 op，但 latch/matprep/IAR 机制才是让它正确且低成本的关键。

如果你读过 LLVM NVPTX 后端，和 latch 最接近的类比是 wgmma accumulator group 之前隐式发生的寄存器文件写入，而和 IAR 最接近的类比是 base+index 寻址模式；不同的是 TPU 把索引寄存器公开为架构对象（IAR0/IAR1），由 SetIar op 显式加载，并由 VectorLoadIndexed op 隐式消费。这里没有逐元素 index 操作数；IAR 就是 index。latch 的 weight-stationary 方式与 wgmma 的 accumulator-stationary 类似：一次 latch 会在多次 matmul step 间摊销。

本页由三个单元构成，对应这三个工作，另加把它们串起来的逐代 cost-model 处理。每个单元都会列出构造该 op 的 builder 函数、它写入的精确 LloInstruction 字段偏移，以及逐代字段/计数表。

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

IAR 字段布局：LloInstruction+0x50→+0x6c 处的 64-bit 值，低 32 位 = index，bit 32 = present；由 CreateVectorSetIarHelper 构建，受 IarsPerTensorCore（每代都是 2）限制。Lane/Sublane/Raw 的划分存在于 ISA bundle-slot opcode 中，而不是该值中。
latch-op family（0x8d..0x96）及其 LloInstruction 字段映射：GainLatchMode @+0x40，latch index @+0x42，MSR @+0x44（opcode-multiplexed），unit-id / MXU-quadrant + source-bus 打包进 control word @+0x0b。
逐代 matprep 机制：GL/GF 给每个 matprep variant 一个固定的二分搜索 perf row；VF 将它折叠进 matmul-format 表外加 modifier-keyed reservation；PF 将它折叠进 latch ops；JF 将 gains 的转置折叠进 matmul opcode。
first-latch overrun handshake（GainLatchModeHasOverrunChecks，vtable +0x358），它决定序列中的第一个 latch 是否会被编号；只有 Viperfish 上才非平凡。


Latch op family	`kVectorLatchLsf` `0x8d` … `kVectorLatch3Msk` `0x96`（10 个 opcode）
Matprep op family	`kVectorMatprepSubr` `0x97`/`0x98`，`kVectorMatprepMubr` `0x99`/`0x9a`
IAR op family	`kVectorReadIar` `0x01`，`kVectorSetIarLane`/`Raw`/`Sublane` `0x02`/`0x03`/`0x04`
IAR value field	`QWORD[LloInstruction+0x50] → +0x6c`；bit 32 present，低 32 index
IarsPerTensorCore	`DWORD[Target+0x4a8]` = 所有世代均为 2（`IAR0`/`IAR1`）
Latch builder	`LloInstruction::CreateVectorLatchLsf` `sub_1D4D7AA0`，`CreateVectorLatchHelper` `sub_1D4D8360`
IAR setter	`LloInstruction::CreateVectorSetIarHelper` `sub_1D4DF080`
Latch sequencing	`MxuAssigner::SetLatchIndices` `sub_10F3B4C0`，`LatchLhs` `sub_10F3B5E0`
Sequence record	`MxuSequence`，sizeof `0x78`，五个 `{ptr,count,cap}` 列表
Home slot	`VectorExtended`（与 matmul 共享）；matres → `VectorResult`

IAR — Index Address Register

用途

IAR 是 TensorCore 的 gather-index 寄存器。VectorLoadIndexed op 读取 base + IAR * stride 处的操作数池，但它不携带显式 index 字段；IAR 寄存器就是逐元素 gather index。使用时，SetIar op 将一个 index 值加载到两个架构寄存器之一（IAR0 / IAR1）；后续 VectorLoadIndexed0 / Indexed1 消费所选寄存器；ReadIar0 / ReadIar1 将其排回 vector register。这是 SparseCore/TensorCore embedding-gather 路径下面的寻址原语。它不同于 SparseCore TEC TileSpmemLoadIndexed，后者确实携带显式逐元素 Index VREG。

IAR Value Field Layout

IAR 值位于 LloInstruction 的 "modifier" 子对象 +0x50（同一个子对象也在 +0x68 保存 precision_type）。iar() 将它作为单个 64-bit 字段读取：

function LloInstruction::iar(this):       // sub_1D4E7120
    sub = *(QWORD*)(this + 0x50);          // the modifier sub-object
    if (sub == nullptr) return 0;          // no IAR present
    return *(QWORD*)(sub + 0x6c);          // 64-bit packed field
```text

builder 是 `CreateVectorSetIarHelper`。它把 index 写入低 32 位，并在同一个 qword 的 byte 4 写入 present-byte，从而设置 bit 32：

```c
function CreateVectorSetIarHelper(opcode, iar_value, source, region):   // sub_1D4DF080
    // bound the index against the per-TensorCore register count
    check(iar_value < Target::IarsPerTensorCore())    // sub_1D617280 → DWORD[Target+0x4a8]
    check(opcode_produced_register_type[source.opcode] == 4)  // source must be a VREG producer
    op = LloInstruction::New(opcode, {source}, region)
    sub = op + 0x50                       // allocate the 0xf0 modifier sub-object if absent (zeroed)
    *(DWORD*)(sub + 0x6c) = iar_value      // sub_1D4DF169 — low 32 bits = the index value
    *(BYTE *)(sub + 0x70) = 1              // sub_1D4DF16C — byte 4 of the qword ⇒ BIT 32 = present
    return op

所以 iar() == (1u64 << 32) | iar_value。 "IAR present" 位是 bit 32；index 值是低 32 位。CreateVectorSetIarRaw（sub_1D4DF260）只是 CreateVectorSetIarHelper(opcode = 0x03, …)；SetIarLane/SetIarSublane/SetIarRaw 全部路由到同一个 helper。

QUIRK — Lane / Sublane / Raw 的区别不在 64-bit 值中。三种形式的 value field 完全相同。硬件写入模式编码在 ISA bundle-slot opcode 中（见下文）；LLO-opcode 编号（0x02/0x03/0x04）和 bundle-slot opcode（2/3/4）并不对齐：LLO SetIarRaw = 0x03 映射到 slot opcode 4，而 LLO SetIarSublane = 0x04 映射到 slot opcode 3。若重新实现时假设 LLO-opcode == slot-opcode，就会把 Sublane 与 Raw 编码交叉弄错。

Bundle-Slot Encoding (PxC / TensorCoreVectorStore)

在 ISA 层面，三个 SetIar 形式只在 word@0x18 的一个 5-bit opcode 子字段中不同；它们的 operand accessor 按字节完全相同。IAR register select 是单个 bit，这也把 IarsPerTensorCore 固定为 2。

NOTE — bit 编号。 本页所有绝对 bit 位置都是 LSB-first，与 Bundle Model 中记录的通用 v5+ packer 约定一致：bit 0 是 byte 0 的最低有效位，因此 word@0x18 bit 13 是 byte 0x18 处 8-byte little-endian word 的 bit 13，而 predicate-mask word@0x18 & 0x3e0000000 选择同一个 word 的五个 bit 33..37。encode/decode 路径中没有任何 MSB-first 顺序。

形式	匹配 predicate (`word@0x18 & 0x3e0000000`)	slot opcode	accessor `sub_ADDR`
`SetIarLane`	`== 0x40000000`	2	`sub_1EE390E0`
`SetIarSublane`	`== 0x60000000`	3	`sub_1EE39100`
`SetIarRaw`	`== 0x80000000`	4	`sub_1EE39120`

字段	Bit 位置	宽度	已确认 accessor
`IarField`（哪个 IAR 寄存器）	`word@0x18` bit 13	1 bit（`IAR0`/`IAR1`）	`sub_1EE3B380`（`>>13 & 1`）
`VsrcField`（source VREG）	`byte@0x1b`	5 bits	`sub_1EE3B360`（`& 0x1f`）

read/use 侧位于 TensorCoreVectorLoad family（major opcode word@0x10 bits[60:62]=7）中，word@0x18 bit 11 处的 2-bit 子字段选择形式：VectorLoadIndexed0 = 0，VectorLoadIndexed1 = 1，ReadIar0 = 2（ReadIar1 是 complement form，这里未 bit-resolved 到单个整数）。indexed-load 携带 DestVreg（6-bit @bit5）、Stride（4-bit @byte0x17）、跨 word 的 2-bit BaseAddress，以及 4-bit SublaneMask，但没有 Index 字段。

IarsPerTensorCore — 寄存器数量

IarsPerTensorCore() 是单条指令的 accessor，且其值不是代码常量：

function Target::IarsPerTensorCore(this):   // sub_1D617280
    return *(uint32*)(this + 0x4a8);
```text

`Target+0x4a8` 的唯一写入者是共享的 `Target::Init`（`sub_1D60FC20`）；没有逐代 `*Target` 构造函数会写它。`Init` 从嵌入的逐代 `*_chip_parts.binarypb` proto 的 `VectorIsa` 字段 7 读取它（运行时通过 `embed://tpu_chip_parts/<version>_chip_parts.binarypb` 加载）。从嵌入 blob 抽取出的值在**每一代都是 2**，IAR 文件因此跨代稳定，并与 1-bit `IarField` 上限匹配。

| Gen（codename） | IarsPerTensorCore（`Target+0x4a8`） | 来源 |
|---|---|---|
| v2 Jellyfish / v3 Dragonfish | 2 | `jellyfish`/`dragonfish` blob `VectorIsa.f7` |
| v4 Pufferfish | 2 | `pufferfish` blob |
| v5p Viperfish | 2 | `viperfish` blob |
| v6e Ghostlite | 2 | `ghostlite` blob |
| v7x (6acc60406) | 2 | `6acc60406` blob |

> **NOTE —** 这个值是数据，不是代码，因此上面的版本 pin 必要但不充分：未来某个芯片的 `VectorIsa.f7` 若不同，就会拓宽 `IarField`。1-bit `IarField` accessor（`sub_1EE3B380`）和计数 2 只对这个 binary 的 chip-parts blob 相互一致。加载路径见 [Chip-Parts Binarypb](../targets/chip-parts-binarypb.md)。

### IAR Cost — Perf-Row Sentinel 划分

cost classifier 以 sentinel `S = ((iar & 0x1ffffffff) == 0x100000000)` 为键选择 indexed-memory perf row；也就是 *IAR present 且 index 值为零*（对齐 base gather，较便宜的 row）与非零 index 值（offset-add row）的区别。Ghostlite classifier `GetGhostliteInstruction`（`sub_1C8B1740`）的七个 IAR-class arm 按字节精确如下；四个 `ReadIar`/`SetIar` arm 还会在 bit 32 未设置时额外 `FATAL`（`"iar.has_value()"`），而三个 indexed-memory arm（`LoadIndexed`、`StoreIndexed`、`StoreIndexedMsk`）会无条件读取 `iar()`，并容忍缺失的 IAR。

```c
// the IAR-class arms of GetGhostliteInstruction, verified in the decompile
case 0x01 ReadIar:        if (!(iar & 1<<32)) FATAL;  return 2*((u32)iar != 0) + 0x18c;  // 0x18c / 0x18e
case 0x02 SetIarLane:     if (!(iar & 1<<32)) FATAL;  return 0x1d5 - ((u32)iar == 0);     // 0x1d4 / 0x1d5
case 0x03 SetIarRaw:      if (!(iar & 1<<32)) FATAL;  return 0x1d9 - ((u32)iar == 0);     // 0x1d8 / 0x1d9
case 0x04 SetIarSublane:  if (!(iar & 1<<32)) FATAL;  return 0x1d7 - ((u32)iar == 0);     // 0x1d6 / 0x1d7
case 0x32 LoadIndexed:    /* no bit-32 gate */         return 2*((iar & 0x1ffffffff) != S) + 0x188; // 0x188 / 0x18a
case 0x40 StoreIndexed:   /* no bit-32 gate */         return ((iar & 0x1ffffffff) == S) ^ 0x1d1;    // 0x1d0 / 0x1d1
case 0x44 StoreIndexedMsk:/* no bit-32 gate */         return ((iar & 0x1ffffffff) == S) ^ 0x1d3;    // 0x1d2 / 0x1d3

LLO op	名称	arm `sub_ADDR`	sentinel-S row	non-S row
`0x01`	`kVectorReadIar`	`sub_1C8B1A94`	`0x18c`（low32==0）	`0x18e`
`0x02`	`kVectorSetIarLane`	`sub_1C8B19DE`	`0x1d4`	`0x1d5`
`0x03`	`kVectorSetIarRaw`	`sub_1C8B1A71`	`0x1d8`	`0x1d9`
`0x04`	`kVectorSetIarSublane`	`sub_1C8B1959`	`0x1d6`	`0x1d7`
`0x32`	`kVectorLoadIndexed`	`sub_1C8B1926`	`0x188`	`0x18a`
`0x40`	`kVectorStoreIndexed`	`sub_1C8B19AD`	`0x1d0`	`0x1d1`
`0x44`	`kVectorStoreIndexedMasked`	`sub_1C8B197C`	`0x1d2`	`0x1d3`

这些 GhPerf::Instruction ordinal 的符号化 enumerator names 不在 binary 中（没有 ToString）；ordinal 按字节精确，但未命名（这是与 matmul perf row 共享的统一空缺）。

Latch — 加载增益矩阵

用途

latch op 将准备好的静态增益（权重）矩阵加载到 MXU 的逐象限 weight bank 中。加载是 weight-stationary 的：一次 latch 会在多次 matmul step 间摊销。增益以何种方式加载（transpose、dtype packing、byte-plane staging）由 GainLatchMode（GLM）操作数决定。这个 family 有十个 opcode：一对 "load-stationary-from-FIFO"（Lsf），一对普通形式，以及三对 indexed sub-bank 形式。

Op Family 和 Builder

每个 LloRegionBuilder::Vlatch* wrapper 都路由到一个 LloInstruction::Create* constructor，后者再路由到 CreateVectorLatchLsf（LSF 特例）或 CreateVectorLatchHelper（通用）。

`LloOpcode`	名称	builder	constructor
`0x8d`	`kVectorLatchLsf`	`VlatchLsf`	`CreateVectorLatchLsf`（`sub_1D4D7AA0`）
`0x8e`	`kVectorLatchLsfMsk`	`VlatchLsfMsk`	`CreateVectorLatchLsfMasked`（`sub_1D4D8140`）
`0x8f`	`kVectorLatch`	`Vlatch`	`CreateVectorLatch`（`sub_1D4D8900`）
`0x90`	`kVectorLatchMsk`	`VlatchMsk`	`CreateVectorLatchMasked`（`sub_1D4D8C40`）
`0x91`	`kVectorLatch1`	`Vlatch1`	`CreateVectorLatch1`（`sub_1D4D8940`）
`0x92`	`kVectorLatch1Msk`	`Vlatch1Msk`	`CreateVectorLatch1Masked`（`sub_1D4D8C80`）
`0x93`	`kVectorLatch2`	`Vlatch2`	`CreateVectorLatch2`（`sub_1D4D8A80`）
`0x94`	`kVectorLatch2Msk`	`Vlatch2Msk`	`CreateVectorLatch2Masked`（`sub_1D4D8CC0`）
`0x95`	`kVectorLatch3`	`Vlatch3`	`CreateVectorLatch3`（`sub_1D4D8B60`）
`0x96`	`kVectorLatch3Msk`	`Vlatch3Msk`	`CreateVectorLatch3Masked`（`sub_1D4D8D00`）

VlatchI(value, long idx, glm)（sub_1D574580）把它的 long idx 分派到 Vlatch1（0x91）或 Vlatch2（0x93）；indexed latch 选择一个 sub-bank。通用 Create* 路径都经过 CreateVectorLatchHelper。

LloInstruction Field Layout

构造出的 latch op 在以下字段中携带操作数；这些字段由 setter 及其对称 reader 按字节精确恢复：

偏移	字段	Setter / reader	含义
`WORD[+0x00]`	`LloOpcode`	`New()`	`0x8d..0x96`
`BYTE[+0x0a]`	`register_number`	`set_register_number` / `sub_1D5A8E20`	gain-source VREG 编号
`WORD[+0x0b]`	control word	`set_unit_id` / `ValidateAndSetMxuAndSourceBus`	unit-id + source-bus（见下）
`BYTE[+0x40]`	`latch_mode`（GLM）	`set_latch_mode` `sub_1D4D7C20` / `latch_mode` `sub_1D4E7500`	GainLatchMode
`WORD[+0x42]`	`latch_index_in_sequence`	`set_latch_index_in_sequence` `sub_1D4E7960`	由 `SetLatchIndices` 分配
`BYTE[+0x44]`	`matrix_staging_register`（Msr）	`set_matrix_staging_register` `sub_1D4D7D40`	latch-bank / MSR destination

control word WORD[+0x0b] 是两个打包 bitfield：

// LloValue::set_unit_id (sub_12698C00) — the GMR / MXU-quadrant pack
WORD[v+0x0b] = (WORD[v+0x0b] & 0xf8ff) + ((unit & 3) << 8) + 0x400;   // check unit <= 3
//   bits 8-9   : unit_id = which MXU quadrant (0..3) the gain matrix latches into
//   bit  10    : has-mxu flag (0x400)

// source-bus pack (ValidateAndSetMxuAndSourceBus, sub_1D4D7E80)
WORD[v+0x0b] = (WORD[v+0x0b] & 0xc7ff) + ((bus & 3) << 11) + 0x2000;
//   bits 11-12 : VEX source-bus (0..3)
//   bit  13    : has-source-bus flag (0x2000)
```text

> **GOTCHA —** `WORD[op+0x42]` 是 latch family（`0x8d..0x96`）的 latch index，但 `BYTE[op+0x42]` 是 *load-LMR* family（`0xaa`/`0xab`）的 MSR。它们共享字节地址，但适用于**互不相交**的 opcode family，因此单个 op 内不存在 aliasing。若重新实现时不先检查 opcode family 就读取 `+0x42`，会误解码其中一类。

MSR setter 是 opcode-multiplexed；同一个字段名会落在四个不同偏移：

```c
function set_matrix_staging_register(op, msr):   // sub_1D4D7D40
    switch opcode_family(op):
        case 0x9b..0xa5 (matmul):        BYTE[op+0x46] = msr
        case 0x8d..0x96 (latch):         BYTE[op+0x44] = msr   // ← the latch family
        case 0xaa/0xab  (load-LMR):      BYTE[op+0x42] = msr
        case 0xa8       (done-with-gains): BYTE[op+0x41] = msr
        default: FATAL "msr unsupported for opcode"

CreateVectorLatchLsf — LSF 构建序列

CreateVectorLatchLsf 是规范的 latch constructor（VlatchLsf 是 append 它的 wrapper）。它检查 gain source 和 GLM，然后写入字段：

function CreateVectorLatchLsf(gain_src, glm, unit_id, region):   // sub_1D4D7AA0
    if (opcode_produced_register_type[gain_src.opcode] != 4)     // gain source must be reg-type 4
        UpdateStatus("chunk->ProducesVreg()")                    // slow diagnostic path otherwise
    if (glm > 0x33 || !bittest(0xf0000003c0c03, glm))            // LSF GLM-validity mask
        FATAL "LSF latch mode not expected."
    op = LloInstruction::New(0x8d /*kVectorLatchLsf*/, {gain_src}, region)
    set_latch_mode(op, glm)                                      // BYTE[op+0x40]
    set_matrix_staging_register(op, 1)                           // BYTE[op+0x44] = 1 (LSF staging slot)
    ValidateAndSetMxuAndSourceBus(unit_id, op)                   // WORD[op+0x0b] unit-id (+ src-bus)
    return op
```text

`VprepareForLatch`（`sub_1D573BA0`）在 constructor 前运行：如果某代并不原生支持 `SupportsGainLatchMode(glm)`（vtable `+0x368`），它会在重新检查前把 gain source 改写成软件 byte-plane 表示。两个 constructor 接受不同的 GLM 集：

| Constructor | GLM-validity mask | 接受 GLM |
|---|---|---|
| `CreateVectorLatchLsf` | `0xf0000003c0c03` | `{0,1,10,11,18,19,20,21,48,49,50,51}`（bf16、F8E5M2、S8、fp8-conv） |
| `CreateVectorLatchHelper` | `0xf000003fffc3f` | `{0-5,10-25,48-51}`（完整集合，含 F8E4M3FN/F32 和 nibble fmt7/8） |

`ValidateAndSetMxuAndSourceBus`（`sub_1D4D7E80`）限制 MXU id（`>= 0`，`< MxusPerTensorCore()` = `Target+0x4ac`），写入 unit-id，并且仅当 `HasVexSourceBuses()`（vtable `+0x408`，**仅 Pufferfish 为 true**）且 `LloOpcodeUsesSourceBus(op)`（`0x8f..0x96` 为 true，`0x8d`/`0x8e` LSF 形式为 false）时写入 source-bus。因此 VEX source-bus 字段只在 Pufferfish（v4）上填充，并且只用于非 LSF latch op。

### First-Latch Overrun Handshake

`SetLatchIndices` 为序列中的每个 latch op 分配程序顺序 index，但**第一个** latch 只有在其 GLM 携带 overrun checks 时才会编号。该 gate 是逐代 `GainLatchModeHasOverrunChecks`（vtable `+0x358`）：

```c
function SetLatchIndices(span<MxuSequence*>):    // sub_10F3B4C0
    for each seq in span:
        for idx = 0 .. seq.latches.count - 1:        // latches list @ seq+0x18, count @ seq+0x20
            op = seq.latches[idx]
            check LloOpcodeIsVectorLatch(op)         // (opcode - 0x8d) < 0xa, else FATAL
            tgt = op.region.module.target            // [[op+0x10]+0x38]+0x10
            glm = latch_mode(op)                     // BYTE[op+0x40]
            has_overrun = tgt.vtbl[+0x358](glm)      // GainLatchModeHasOverrunChecks
            if (idx == 0 && !has_overrun): break      // first latch, no overrun ⇒ abandon sequence
            set_latch_index_in_sequence(op, idx)      // WORD[op+0x42] = idx

五个世代中有四个是平坦 FALSE；它们的第一个 latch 永远不会编号。Viperfish 是唯一具有该 handshake 的世代，并且只针对宽的非 bf16 NO_XPOSE mode：

function ViperfishTarget::GainLatchModeHasOverrunChecks(glm):   // sub_1D49AB20
    if (LatchModeIsTranspose(glm)) return false;                // transpose ⇒ no overrun
    fmt = GainLatchModeToMatmulDataFormat(glm);
    return MatmulDataFormatIsIntegral(fmt) | ((fmt - 3) < 2);   // ⇒ fmt ∈ {3,4,5,6,7,8}
```text

| Gen | `GainLatchModeHasOverrunChecks`（`+0x358`） | `HasMsrOverrunChecks` |
|---|---|---|
| Jellyfish（`sub_1D4925E0`） | `FALSE`（always） | `FALSE` |
| Dragonfish（`sub_1D4901C0`） | `FALSE` | `FALSE` |
| Pufferfish（`sub_1D494880`） | `FALSE` | `FALSE` |
| Viperfish（`sub_1D49AB20`） | non-transpose AND fmt∈{3..8} → GLM `{14,16,18,20,22,24}` | **`TRUE`**（`sub_1D49AAC0`） |
| Ghostlite（`sub_1D497940`） | `FALSE` | `FALSE` |
| base Target（`sub_1D61D8C0`） | `LogFatal` stub | `LogFatal` |

> **NOTE —** Viperfish（TPU v5p）是唯一在 gen 级别拥有 MSR/first-latch overrun handshake 的世代，这正是它的逐 GLM `+0x358` override 是唯一非平凡主体、且 overrun-cost reservation 位于 Viperfish namespace 中的原因。完整 overrun 行为（first-latch index assignment、MSR reservation cost）见 [Latch Assignment & Overrun](../sched/latch-assignment-overrun.md)。

---

## Matprep — 逐代暂存操作数

### 用途

Matprep 暂存移动操作数（activations）并为 latching 准备增益矩阵。*真正的* matprep opcode 是 `kVectorMatprepSubr`（`0x97`/`0x98`，sub-row 形式）和 `kVectorMatprepMubr`（`0x99`/`0x9a`，block-row 形式），再加上 `kVectorMatmulLmr` / `kVectorDoneWithGains` / `kVectorLoadGmr` helper（`0xa5`/`0xa8`/`0xa9`）。它们不同于 gain-LATCH family（`0x8d..0x96`）。

关键的重新实现事实是，matprep **没有统一的 cost 表示**；每一代都以不同方式表达它，而重新实现者必须复现这种分歧：

| Family | Jellyfish/Dragonfish (v2/v3) | Pufferfish (v4) | Viperfish (v5p) | Ghostlite/GF (v6e/v7) |
|---|---|---|---|---|
| matmul | flat cell（LUT collapse 到 5 instrs） | raw 2-bit plane + base | `matmul_data_format → a2d05c0` | `matmul_data_format → a2d05d0` |
| matprep `0x97..0x9a` | 折叠进 matmul（transpose-of-gains） | 经由 Latch ops `0xdc`/`0xe6` | matprep ops **FATAL**；经由 matmul-fmt + modifier reservation | 固定二分搜索 rows |
| transpose accepted | `{B32}` only | `{B32,CompB16,SegB32,SegB16}` | `{B32,CompB16,SegB32,SegB16}` | `{B32,CompB16,CompB8}` |

### GL / GF — 固定二分搜索 Rows

在 Ghostlite/GF 上，matprep opcode 不在 classifier jump table 中（`opcode-1 ≥ 0x96` 会落到 default-FATAL），所以它们通过 258-entry 二分搜索 remap（`@0x4067dc8`）解析为**固定** perf row；每个 matprep variant 一个 row，*不会*按 data format 展开：

| LLO op | 名称 | `GhPerf::Instruction` | GF flat latency |
|---|---|---|---|
| `0x97` | `kVectorMatprepSubr` | `0x120` | 1 |
| `0x98` | `kVectorMatprepSubrMsk` | `0x121` | (matprep) |
| `0x99` | `kVectorMatprepMubr` | `0x11c` | 1 |
| `0x9a` | `kVectorMatprepMubrMsk` | `0x11d` | 1 |
| `0xa5` | `kVectorMatmulLmr` | `0x154` | −1 default（grid-priced） |
| `0xa8`/`0xa9` | `kVectorDoneWithGains` / `kVectorLoadGmr` | `0x157`（shared） | −1 default |

matprep band `0x11c..0x121` 位于 matmul band `0x124..` 正下方；这些 row 携带 flat latency 1，并通过 resource grid 进行 throughput 计价。flat-latency-1 值来自 GF perf constructor `sub_1C8D3740`。

### VF — 折叠进 Matmul-Format 表 + Modifier Reservation

在 Viperfish 上，matprep opcode `0x97..0x9a` 会在 `GetViperfishInstruction`（`sub_1C8A3300`，default arm `sub_1C8A3E6A`）中 **FATAL**。matmul opcode `0x9b` 读取 `matmul_data_format()` 并索引新的 VF 表 `a2d05c0`：

| `MatmulDataFormat` | dtype | VFinstr ordinal | flat latency | grid（r2 prep / r3 throughput） |
|---|---|---|---|---|
| 1 | f32 | `0xd4` | 131 | r2:7 r3:8 |
| 2 | bf16 | `0xda` | 131 | r2:7 r3:16 |
| 3 | f8e5m2→bf16 | `0xf8` | 131 | r2:7 r3:32 |
| 4 | f8e4m3b11→bf16 | `0xfe` | 131 | r2:7 r3:32 |
| 5 | u8 | `0xe0` | 121 | r3:16 |
| 6 | s8 | `0xe6` | 121 | r3:16 |
| 7 | u4 | `0xec` | 121 | r3:16 |
| 8 | s4 | `0xf2` | 121 | r3:16 |

throughput port `r3` 是逐格式 reservation width：f32=8、bf16=16、fp8=32、int8/int4=16。bf16-class（`0xd4..0xfe`）携带独立 prep port（`r2:7`）和 base latency 131；int-class（`0xe0..0xf7`）去掉 `r2`，使用 121。每个 matmul-format ordinal 后面跟着一组 matprep-stage ordinal（例如 `0xd4` 和 `0xda` 之间的 `0xd5/0xd6/0xd8/0xd9`），它们添加 4-stage systolic-feed pipeline `r4:4 r5:12 r6:20 r7:28`。

matprep stage 不携带独立 classifier ordinal；它们由 `MxuLatencyTable::GetResourceUsage`（`sub_1C8AE5C0`）生成，该函数构建 `MatpushModifier { MatmulDataFormat, is_transpose, Msr }` key，并在 `FlatHashMap<Modifier, array<int,19>>` reservation table 中查找（matprep `r4..r7` stage 是 19 个 `MxuResource` port 中的 4 个）。见 [Matmul-Mode Modifiers](../cost/matmul-mode-modifiers.md)。

### PF — 折叠进 Latch Ops

在 Pufferfish 上，matprep opcode `0x97..0x9a` 同样 FATAL（default arm `sub_1C8A2A08`）。PF 通过单 ordinal 的 Latch / LatchMsk arm 表达 matprep，并由 `SupportsGainLatchMode`（vtable `+0x368`）gate：

| LLO op | 名称 | arm `sub_ADDR` | PFinstr |
|---|---|---|---|
| `0x8f` | `kVectorLatch` | `sub_1C8A2781` | `0xdc` |
| `0x90` | `kVectorLatchMsk` | `sub_1C8A226F` | `0xe6` |

`0xdc`/`0xe6` 是 PF "MXU matprep band" 的入口；PF 的 matprep 就是这一对单独的 Latch/LatchMsk，通过 resource-grid port 进行 throughput 计价。

### JF / DF — Transpose-of-Gains 折叠进 Matmul

在 Jellyfish/Dragonfish 上没有独立的 matprep classifier；`CycleTableInstruction`（`sub_1C89CA80`）通过 LUT 将 11 个 `MatmulDataFormat` 值折叠为 5 条 instruction，而 gain 的转置折叠进 matmul opcode：`EmitVectorMatmul`（`sub_140B92C0`）基于 `DoneWithGainsMode == 2 (TRANSPOSED)` 分派 VEOpcode：`0x9b → 0/4`、`0x9d → 2/6`、`0x9e → 1/5`。独立 transpose 只接受 `VxposeMode 0`（B32）→ VEOpcode `0xf`；其他所有 mode 都 FATAL（"JFC/DFC only support B32 transpose instructions"）。

> **QUIRK —** matprep 的*表示*在各代之间迁移：JF 将它吸收到 matmul 中，PF 吸收到 Latch ops 中，VF 吸收到 matmul-format 表外加 modifier reservation 中，只有 GL/GF 给每个 matprep variant 一个专用固定 perf row。若重新实现时假设各代共享同一种 matprep cost model，会错误计价五代中的四代。

---

## MxuSequence Record

### 用途

`MxuSequence` 是 `MxuAssigner` 迭代的逐序列记录：`SetLatchIndices` 对 latch 排序，`LatchLhs` 划分增益矩阵并发出 latch+matmul+matres ops，而 `AllocateMrb`/`Bounce`（输出侧）分配 result-FIFO 地址和 MSR bank。它持有五个 instruction list；逐 instruction 的 "latch state" 分布在成员 `LloInstruction` 上，而不是存成平坦 scalar。

### 布局（sizeof `0x78`）

从 deleter `default_delete<MxuSequence>::operator()`（`sub_14504C00`）恢复；它释放五个 `{ptr, count, cap}` 列表，然后 `free(seq, 0x78)`：

| 偏移 | 列表 | 元素 opcode / consumer |
|---|---|---|
| `+0x00` | list0（setup/head latches） | head-of-sequence（按类别推断） |
| `+0x18` | latches / matpushes | `0x8d..0x96` — `SetLatchIndices` count @`+0x20`；`Bounce` MSR stamp |
| `+0x30` | list2（prep / xpose aux） | matprep/transpose `0xa6`/`0xa7`（按类别推断） |
| `+0x48` | matreses | `0x152` — `AllocateMrb` pop；`LatchLhs` ΣPackingFactor，count @`+0x50` |
| `+0x60` | matmuls | `0x9b`/`0xa3` — `AllocateMrb` push；`LatchLhs` balance，count @`+0x68` |
| `0x78` | sizeof | `free(seq, 0x78)` |

`+0x18`、`+0x48`、`+0x60` list 身份按字节精确（由 deleter 和三个独立 consumer 确认）；`+0x00` 和 `+0x30` 身份按类别推断。完整 record 和 `set_mxu` commit 位于 [MxuSequence / SequenceInfo](../sched/mxu-sequence-struct.md)。

### LatchLhs — 增益矩阵划分

`LatchLhs`（`sub_10F3B5E0`）是 latch+matmul+matres ops 的生产者，`SetLatchIndices` 稍后会给这些 op 编号。它按 transpose op 对 LHS 分组，运行逐 MXU capacity guard，然后用每个 op 的 MXU quadrant 标记重建序列：

```c
function LatchLhs(target, lhs_span, sequences):   // sub_10F3B5E0
    xpose = BuildXposeSequences(lhs_span)          // vec1 = {0xa6,0xa7}, vec2 = {0x154}
    // capacity guard per sequence
    acc = Σ over matreses of MatmulDataFormatPackingFactor(matmul_data_format(op))
    check( ChunksPerTile() * num_mxus >= acc )     // ChunksPerTile = hwcfg[+0x198]/hwcfg[+0x1a0]
    check( acc % ChunksPerTile() == 0 )            // tile-aligned, num_mxus = Target+0x4ac
    // rebuild per quadrant
    for each matmul:
        q   = program_order & 3                    // the MXU quadrant (0..3)
        glm = GLM_byte_table[matmul_op - 0x9b]     // @0xac0913e: {0×8, 0xb, 0xb} ⇒ plain→0, packed→0xb
        VlatchLsf(value, glm, 0)                    // emit kVectorLatchLsf (sub_1D573EC0)
        WORD[emit+0x0b] = (WORD[+0x0b] & 0xf8ff) | ((q<<8)+0x400)   // set_unit_id(q)
        repeat Vmatmul / Vmatres PackingFactor(fmt)× (K-tile split), each unit_id-stamped

MatmulDataFormatPackingFactor（sub_1D629300）索引表 @0xb53c6bc = {1,2,4,4,4,4,8,8,4,4}（fmt 1..10）；这是驱动 K-tile loop count 的 column-pack factor。unit_id（= MXU quadrant）是 gain-matrix-register bank；GLM 是 latch mode；MSR（输出侧）是 staging bank。

名称	关系
MXU Slot	消费已 latch 的 gains 和 matprep 后的 operand；matmul step 本身
Jellyfish 41-Byte Bundle	v3 `VectorExtended` 编码，用于序列化这些字段
Latch Assignment & Overrun	`SetLatchIndices` + 逐代 overrun handshake（调度侧）
MxuSequence / SequenceInfo	完整 sequence record 和 `set_mxu` commit
Matmul-Mode Modifiers	VF `Modifier → array<19>` matprep reservation table

交叉引用

MXU Slot — 这些子槽喂给的 systolic-array op family；matmul / matprep / latch 共享 VectorExtended
Jellyfish 41-Byte Bundle — 将 latch/matprep 字段编码进 VectorExtended slot 的 v3 bundle
Latch Assignment & Overrun — SetLatchIndices、LatchLhs 和 Viperfish first-latch overrun handshake
MxuSequence / SequenceInfo — 按字节精确的 MxuSequence record 和逐 instruction latch state
MXU Latency Overview — 为 matprep/latch perf row 计价的逐代 reservation matrix
IARs Per TensorCore — Target+0x4a8 register-count 字段及其 chip-parts 来源
Matmul-Mode Modifiers — VF MatpushModifier-keyed array<int,19> matprep reservation
Chip-Parts Binarypb — 提供 IarsPerTensorCore 的嵌入式逐代 proto

Matprep、IAR 和 Latch 子槽 ​

摘要 ​

IAR — Index Address Register ​

用途 ​

IAR Value Field Layout ​

Bundle-Slot Encoding (PxC / TensorCoreVectorStore) ​

IarsPerTensorCore — 寄存器数量 ​

Latch — 加载增益矩阵 ​

用途 ​

Op Family 和 Builder ​

LloInstruction Field Layout ​

CreateVectorLatchLsf — LSF 构建序列 ​

相关组件 ​

交叉引用 ​