MXU 槽

地址适用于 libtpu-0.0.40-cp314 wheel 中的 libtpu.so（BuildID md5 89edbbe81c5b328a958fe628a9f2207d，未 strip）。其他版本会有所不同。

摘要

MXU 槽是驱动 systolic matrix-multiply array 的 bundle 字段，也是 TensorCore ISA 中最重要的单一槽，因为它是通向 TPU 标志性 dense-linear-algebra 吞吐量的唯一路径。它不是一个 opcode，而是贯穿一个或两个物理 bundle 槽的小型指令家族：latch（将 stationary weight matrix 加载进阵列）、matpush（暂存 moving operand）、matmul（推进一个 systolic step）和 matres / result-pop（排空 accumulator）。跨五代硬件，这个家族以五种不同方式编码，但底层契约保持不变：weight-stationary systolic multiply，由共享 vector-register operand pool 供给，并通过独立 result 槽排出。

该槽有两个结构上不同的 encoder 谱系。在 Jellyfish (v2) 和 Dragonfish (v3) 上，MXU 是一个 VectorExtended 槽，其 6-bit opcode 通过 shl/and/or 算术打包进 scratch struct 的 qword 0，并由两级 jump table 解码。从 Pufferfish (v4) 开始，encoder 变成通用 TensorCoreCodecBase 模板：两个 VectorExtended 槽（每个物理 MXU 一个控制区域），每个字段由通用 BitCopy(dst, dst_bit, src, 0, width) 调用放置，并由线性 Opcode::Matches 扫描解码。两个 MXU 控制区域是固定的 −N-bit twin（PF/VF 为 −20，Ghostlite 为 −21，6acc60406 为 −25），打包在一个共享的 8×6-bit operand pool 之上。本页覆盖该槽跨所有代的语义；绝对位位置位于各代 bundle 页面。

如果你读过 LLVM NVPTX 后端，最接近的类比是 wgmma/mma.sync：单个 machine op 将 operand registers、accumulation half 和 data format 都编码为字段，并单独排出结果。TPU 的不同之处在于它显式是多指令的：latch、push、multiply 和 pop 是编译器调度的不同 bundle 槽，而不是一个融合 intrinsic；并且它是 weight-stationary，因此 latch 可在许多 matmul steps 之间摊销。

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

MXU op 家族及其到物理 bundle 槽的映射：latch + matprep + matmul 共享 VectorExtended 槽；matres 使用 VectorResult。
VectorExtended 槽的各代字段布局：opcode、data-format / dtype 子判别符、MXU-id、done-gains / transpose / target 控制位，以及 5-bit（v2/v3）/ 4-bit（v5+）predicate。
opcode 编码规则：JF emitter 如何将 LloOpcode × DoneWithGainsMode 映射到 6-bit opcode，以及 v5+ codec 如何将 7/8-bit opcode 字段拆分为 {matmul, PushGains/latch, transpose} 家族并带有逐 dtype 值。
−N-bit twin 几何和共享 operand pool 模型，这使两个 MXU 控制区域可在一个 bundle 中共存。


Op family (LLO)	`vlatch`/`vlatchi`（latch），`vmatprep.subr`/`.mubr`（matpush），`vmatmul`(`.high`/`.low`/`.mubr`/`.msk`)，`vmatres`(`.add`)
JF LloOpcode values	matprep.subr `0x97/0x98`，matprep.mubr `0x99/0x9a`，matmul `0x9b`，matmul.mubr `0x9c`/`0xa0`，matmul.high `0x9d`，matmul.low `0x9e`，matres `0x152`
JF slot home	latch/matprep/matmul → `VectorExtended`（slot_mask `0x080`，proto `+0x50`）；matres → `VectorResult`（`0x100`，`+0x58`）
JF VE field map	opcode @ abs 29..34（6b），mxu-id @ 27..28（2b），predicate @ 35（5b）— `EncodeVectorExtendedInstruction` @ `0x1e869f00`
v5+ MXU slots	两个 `VectorExtended` 控制区域（每个物理 MXU 一个），共享 operand pool 之上的 `−N`-bit twin
Systolic array	128 × 128（JF–VF）/ 256 × 256（Ghostlite、`6acc60406`）weight-stationary；由 8×128 / 4×256 matpush tiles 供给
Empty-slot mark	predicate `kNeverExecute = 31`（`0xB834CFC`）预填进槽

MXU Op 家族

目的

MXU 是 weight-stationary systolic array。完整 dense matmul C = A·B 不是一条指令，而是编译器发出并由 bundle packer 调度的一序列 LLO ops：latch stationary weight tile、push moving-operand tile、让阵列向前 clock，并排出 result FIFO。op 家族就是该序列的词汇表。

Op 名单

Jellyfish op 集直接从 per-op LloOpcodeIsVector* 分类器函数读取，每个都是很小的 (opcode − base) < n 或 mask 测试。分类器值是 CERTAIN，因为它们是二进制中的算术常量，而非推断。

LLO mnemonic	`LloOpcode`	槽	作用
`vmatprep.subr`（`+.msk`）	`0x97`/`0x98`	VectorExtended	push moving operand，sub-row 形式
`vmatprep.mubr`（`+.msk`）	`0x99`/`0x9a`	VectorExtended	push moving operand，block-row 形式
`vmatmul`	`0x9b`	VectorExtended	一个 systolic step
`vmatmul.mubr`（`+.msk`）	`0x9c`/`0xa0`	VectorExtended	conv block-row matmul
`vmatmul.high`	`0x9d`	VectorExtended	high-half accumulator step
`vmatmul.low`	`0x9e`	VectorExtended	low-half accumulator step
`vlatch` / `vlatchi`	（latch path）	VectorExtended	将 stationary weights latch 进阵列
`vmatres`（`+.add`）	`0x152`	VectorResult	排出 result FIFO（`.add` = accumulate）

反编译中验证的分类器算术：

// platforms_deepsea::jellyfish — the per-op LloOpcode classifiers
LloOpcodeIsVectorMatprepSubr(op):  return (uint16)(op - 151) < 2;   // {0x97, 0x98}   @ 0x1d60c400
LloOpcodeIsVectorMatprepMubr(op):  return (uint16)(op - 153) < 2;   // {0x99, 0x9a}   @ 0x1d60c3e0
LloOpcodeIsVectorMatmulMubr(op):   return ((op - 156) & 0xFFFB)==0; // {0x9c, 0xa0}   @ 0x1d60c3c0
// matmul=0x9b, matmul.high=0x9d, matmul.low=0x9e read from the EmitVectorMatmul dispatch;
// matres=0x152 from EmitVectorMatres @ 0x140b9600 (cmp 0x152).
```text

相同的矩阵 ops 也注册为 MLIR LLO dialect operations — `VectorMatmulOp`、`VectorMatmulMubrOp`、`VectorMatprepSubrOp`、`VectorMatprepMubrOp`、`VectorMatresOp`、`VectorLatchOp`、`VectorLatchIOp`、`VectorDoneWithGainsOp`、`VectorMoveEvenAccLowOp` — 确认 dialect 表面与编码 opcodes 匹配。

> **注意 —** 连续的 `0x97..0xa0` 块包含整个 matprep/matmul 家族，但 `matres` 位于很远的 `0x152`。这两个区域反映两个物理槽：matprep/matmul/latch 供给阵列（`VectorExtended`）；matres 排出它（`VectorResult`）。假设单一连续 MXU opcode 块的重新实现会误分类 result-pop。

### Operand 角色 — Stationary 与 Moving

阵列有两个 operand，由两个不同 ops 供给：

- **Stationary（weights / gains）** — 由 **latch** ops（`vlatch`/`vlatchi`，v5+ 的 `LoadMatrixRegister*` / `PushGains*` 家族）加载。stationary operand 会跨许多 matmul steps 常驻阵列，直到重新 latch。它*如何*加载 — transpose、packing、dtype staging — 由 `GainLatchMode` enum（v2/v3）或逐 dtype `PushGains<fmt>` / `Pushmatrix<fmt>` opcode（v5+）表示。
- **Moving（activations / multiplicand）** — 由 **matprep** ops（`vmatprep.subr` = sub-row，`.mubr` = multiplicand-block-row）暂存。moving operand 被暂存进 matrix-staging register，即 `MATPUSH_TARGET_MSRA` 或 `MATPUSH_TARGET_MSRB`（两个字符串都存在于 `.rodata`），随后通过阵列 clock。

matpush tile 粒度由硬件 perf-counter help string 固定，字符串逐字节存在于二进制中：*"This counts the number of 8x128 or 4x256 matrices that are pushed to MXUn by a vmatpush instruction."* 因此每次 push 提供一个 **8×128**（sublanes × lanes）或 **4×256** tile；十六个 8×128 tiles 填满一个 128×128 阵列（JF–VF 几何；Ghostlite/`6acc60406` 将阵列扩宽到 256×256，见 [Systolic-Array 几何](#systolic-array-几何)）。

---

## Jellyfish (v2) — 直接打包的 VectorExtended 槽

### 槽分配

Jellyfish 有**一个** MXU，且每个 bundle 有一个承载 MXU 的槽。MXU ops 没有专用槽；它们复用 vector pipeline 的 `VectorExtended` 和 `VectorResult` 槽。`JellyfishTarget::NumVexSlots()`（`0x1d4912c0`）返回 **1**，正好每个 bundle 一个 `VectorExtended` 槽；`NumVsSlots()`（`0x1d4912a0`）返回 **3**（供给它的三个 vector-source 端口 vs0/vs1/vs2）。

| MXU op | Bundle 槽 | `slot_mask` bit | proto offset |
|---|---|---|---|
| latch / matprep / matmul | `VectorExtended` | `0x080` | `proto+0x50` |
| matres | `VectorResult` | `0x100` | `proto+0x58` |

emitter dispatch：

```text
JellyfishEmitter::EmitVectorMatmul   @ 0x140b92c0  ── matmul (all forms)
JellyfishEmitter::EmitVectorLatch    @ 0x140b8c20  ── latch + matprep
JellyfishEmitter::EmitVectorMatres   @ 0x140b9600  ── matres (op 0x152)
  ├─ EmitVectorExtendedInstruction   @ 0x140b4f80  ── builds the VE proto submessage
  ├─ EmitVectorResultInstruction     @ 0x140b53e0  ── builds the VR proto submessage
  ├─ AddMxuNumToVectorExtended       @ 0x140b8da0  ── stamps the 2-bit MXU id (proto +0x70)
  ├─ AddMxuNumToVectorResult         @ 0x140b9680
  └─ CheckMxuNum                     @ 0x140b9780  ── JF: mxu must == 0; Dragonfish: free

VectorExtended 字段布局

EncoderJf::EncodeVectorExtendedInstruction（0x1e869f00）将字段 OR 进 struct byte 0x0C 处的 8-byte little-endian word（按 12-byte-strip law，即绝对 bundle bit 96）。由于 struct byte 0x0C 是 bit 96，qword-0 shifts 可直接读作其绝对 bundle 位位置。字段位置已从反编译 mask 逐字节验证：

// EncoderJf::EncodeVectorExtendedInstruction @ 0x1e869f00 (verified)
word  = struct[0x0C];
word  = (predicate & 0x1F) << 35 | word & 0xFFFFFF07FFFFFFFF;   // predicate @ abs 35
word  = (mxu_id    & 0x03) << 27 | word & 0xFFFFFFFFE7FFFFFF;   // mxu-id    @ abs 27..28
// opcode field cleared by 0xFFFFFFF81FFFFFFF (bits 29..34); each VEopcode value ORed/added in:
//   opcode 0 → word &= ~mask;            opcode 1 → | 0x20000000;   opcode 2 → + 0x40000000;
//   opcode 3 → + 0x60000000;             ... opcode 0x22 → + 0x540000000   (35-case jump table)
```text

| 字段 | 来源 | abs bits | 宽度 |
|---|---|---|---|
| predicate | `EncodePredication & 0x1F` | 35 | 5 |
| opcode（`VectorExtendedOpcode`） | 35-case jump table | 29..34 | 6 |
| mxu-id（unit） | proto `+0x64 & 3` | 27..28 | 2 |
| operand vregs | proto `+0x6c` 等 | 依 sub-mode 而定 | 各 5 |

opcode clear-mask `0xFFFFFFF81FFFFFFF` 与 decoder（`DecoderJf::DecodeVectorExtendedSlot` @ `0x1e854000`）用于提取 abs 29..34 上 6-bit opcode 的 mask 相同，因此 encode 和 decode 按位一致；这将 opcode / mxu-id / predicate 字段提升到 CERTAIN 级交叉确认。绝对位置上下文见 [Jellyfish 41B Bundle](bundle-jf-41b.md#vector-extended--mxu-struct-0x0c--encodevectorextendedinstruction--0x1e869f00)。

6-bit `VectorExtendedOpcode` 空间分为三个家族（分类器范围来自 `ProtoUtils::IsMatrixMultiply` @ `0x1e875b20`、`IsPushGains` @ `0x1e875b80`、`IsTranspose` @ `0x1e875b40`、`IsRpu` @ `0x1e875b60`）：

| VEopcode range | 分类器 | 家族 |
|---|---|---|
| `0..6` | `IsMatrixMultiply`（op < 7） | matmul（op 3 = 仅 staging；`VectorExtendedUsesData` @ `0x1e876160` 返回 op != 3） |
| `7..12` | `IsPushGains`（7..12） | weight-latch（`GainLatchMode 0..5` 范围） |
| `17,18` | `IsTranspose` | matrix transpose |
| `17..34` | `IsRpu` | reduce / permute / transpose 家族 |

### Matmul Opcode 编码

`EmitVectorMatmul`（`0x140b92c0`）将 matmul `LloOpcode` 转换成小的 `VectorExtendedOpcode`。选择器是 **`DoneWithGainsMode`** 参数，*不是* data format；dispatch 测试 `dwg_mode == kTransposed`（值 2）。反编译 `switch` 明确无歧义：

```c
// JellyfishEmitter::EmitVectorMatmul @ 0x140b92c0 (a2 = LloOpcode, a5 = DoneWithGainsMode)
switch (a2):
  case 0x9b:  veopcode = 4 * (dwg != 2) + 0;  break;   // matmul        → 0 if transposed, else 4
  case 0x9e:  veopcode = 4 * (dwg != 2) + 1;  break;   // matmul.low    → 1 if transposed, else 5
  case 0x9d:  veopcode = 4 * (dwg != 2) + 2;  break;   // matmul.high   → 2 if transposed, else 6
  default:    NoteError("unhandled LLO opcode for matrix multiply: %s");
EmitVectorExtendedInstruction(veopcode, ...);
AddMxuNumToVectorExtended(mxu_num);

`LloOpcode`	`dwg == kTransposed`（2）	else
`vmatmul`（`0x9b`）	VEopcode 0	VEopcode 4
`vmatmul.low`（`0x9e`）	VEopcode 1	VEopcode 5
`vmatmul.high`（`0x9d`）	VEopcode 2	VEopcode 6

注意 — EmitVectorMatmul 中的 cmp $2 选择器测试的是 DoneWithGainsMode（参数 a5），其中值 2 是 DoneWithGainsMode::kTransposed，而不是 MatmulDataFormat。同一函数携带由 "JF/DF have only one GSF"（jellyfish_emitter.cc:1710）保护的断言 dwg_mode != DoneWithGainsMode::kTransposed。MatmulDataFormat 是独立参数，不会在 Jellyfish 上驱动此 opcode dispatch。

Latch Opcode 编码

EmitVectorLatch（0x140b8c20）携带一个限定为 0..5 的 GainLatchMode 参数（if (mode >= 6u) 是 fatal 路径），并索引 VA 0xaef42ac 处的 6-entry .rodata 表以获得 latch VectorExtendedOpcode。表字节直接从 ELF 读取，逐字节精确为 {7, 10, 9, 12, 8, 11}：

// JellyfishEmitter::EmitVectorLatch @ 0x140b8c20 (verified)
if (gain_latch_mode >= 6u) fatal(...);            // GainLatchMode must be 0..5
veopcode = dword_AEF42AC[gain_latch_mode];        // {7,10,9,12,8,11} read from ELF off 0xaef42ac
EmitVectorExtendedInstruction(veopcode, vreg, ...);
```text

| `GainLatchMode` | VEopcode | | `GainLatchMode` | VEopcode |
|---|---|---|---|---|
| 0 | `0x07` | | 3 | `0x0c` |
| 1 | `0x0a` | | 4 | `0x08` |
| 2 | `0x09` | | 5 | `0x0b` |

这六个 VEopcodes 正好落在 `IsPushGains` 范围（7..12）内。完整跨代 `GainLatchMode` enum 丰富得多（`NO_XPOSE_F32`/`HI_F32`/`LOW_F32`、`S4`/`S8`/`U4`/`U8`、soft/nibble/packed/FP8 staging modes，以及对应的 `XPOSE_*` transposed 变体）；v3 只使用 6-entry 表覆盖的小型 `{F32, HI_F32, LOW_F32, S8/U8, …}` 子集。更丰富的 modes 在后续代成为具名 opcode 家族（见下文）。

### VectorResult（matres）字段布局

`EncoderJf::EncodeVectorResultInstruction`（`0x1e865ae0`）共享 struct `0x0C` 处同一个 word；`VectorExtended` 和 `VectorResult` 可在一个 bundle 中共存于不相交位范围。从 encoder mask 解码得到：

| 字段 | 来源 | abs bits | 宽度 |
|---|---|---|---|
| predicate | `& 0x1F` | 22..26 | 5 |
| result type / format | proto `+0x40` | 20..21 | 2 |
| result mode | proto `+0x44` | 18..19 | 2 |
| dest vreg | proto `+0x48` | 依 mode 而定 | 5 |

2-bit result-mode（proto `+0x44`，值 0/1/2）选择从哪个 MRF/MSR FIFO 排出结果，并门控 destination-vreg shift；`matres.add`（独立 LLO opcode）累加进 result accumulator，而不是覆盖。`AddMxuNumToVectorResult`（`0x140b9680`）还会写入 source MXU id。

### 单 MXU 约束

`JellyfishEmitter::CheckMxuNum(int)`（`0x140b9780`）在普通 Jellyfish 上 `CHECK` `mxu == 0`（device == `kJellyfishIdentifiers`，断言 *"0 != mxu"*）。在 **Dragonfish**（`kDragonfishIdentifiers`，v3 multi-MXU 封装）上，2-bit MXU-id 字段允许非零并写入槽中。因此 2-bit MXU-id 字段存在于共享的 JF/DF（v2/v3）codec 中，在 Jellyfish 上固定为 0，在 Dragonfish 上可用；encoder 其他方面相同（`EncoderDf::EncodeVectorExtendedInstruction` @ `0x1e85e520` 共享布局）。

---

## Pufferfish (v4) — 双 MXU Codec 起点

### 槽翻倍与 −20 Twin

Pufferfish 将 Jellyfish 的单个 `VectorExtended` 槽翻倍为**两个**独立 MXU 控制槽，并切换到 `TensorCoreCodecBase` 模板 encoder（无 scratch struct、无 header strip；每个字段都由 `BitCopy(buf, dst_bit, &field, 0, width)` 调用放置，且 `dst_bit` *就是*绝对 bundle 位）。MXU1（abs 63..82）是 MXU0（abs 83..102）的逐位孪生体，偏移精确为 **−20 bits**。

| 字段 | MXU0 abs | MXU1 abs | 宽度 |
|---|---|---|---|
| sub-op | 83 | 63 | 3 |
| mode / mxu-num | 89 | 69 | 2 |
| opcode | 91 | 71 | 7 |
| predicate | 98 | 78 | 5 |

槽图上下文见 [Pufferfish 51B Bundle](bundle-pf-51b.md#vector-extended--mxu-slots--0x1edb0900-mxu0-0x1ee08060-mxu1)。

### Opcode 字段 — Matmul 加宽并携带物理 MXU 编号

abs 91 处的 7-bit opcode 选择 op 家族；对 **matmul**，该字段加宽为 **9 bits**（abs 89..97），低两位 @ 89..90 携带*物理* MXU number 0..3。这是 v4 中 bundle 槽与物理阵列正交性的起点，由 decode-side `Opcode::Matches` masks 确认：

| Mnemonic | opcode 值 | abs bits | 说明 |
|---|---|---|---|
| `Noop` | predication == 0 | 98..102 | empty |
| `MatmulLow Mxu0..3` | 4, 5, 6, 7 | 89..97（9b） | op-hi 1；mxu-num @ 89..90 |
| `MatmulHi Mxu0..3` | 8, 9, 10, 11 | 89..97 | op-hi 2；mxu-num @ 89..90 |
| `PushGains{Rounded,Low,Hi,Packed,Byte}` | `0x20`..`0x24` | 91..97（7b） | weight-latch |
| `PushGains{Low,Hi,Byte}Masked` | `0x31`, `0x32`, `0x34` | 91..97 | `= 0x2N + masked·0x10` |
| `DoneWithGains{Gsfn,Gsft}` | `0x18`, `0x19` | 91..97 | end-of-gains |
| `Transpose` / `PackedTranspose` | `0x40` / `0x48` | 91..97 | systolic transpose op |

因此 `PushGains opcode = 0x20 + {Rounded 0, Low 1, Hi 2, Packed 3, Byte 4} + masked·0x10`，而 `matmul opcode = (op-hi << 2) | mxu-num`。matmul-vs-PushGains 区分由 opcode-high 值决定。

> **易错点 — 第二个 MXU 由 opcode 字段选择，而不是 bundle 槽，尽管存在两个 MXU *slots*。** Pufferfish 有两个 MXU *control* slots（MXU0/MXU1，即 −20 twin）和**四个**物理 MXU arrays（`mxu_count = 4`；perf counters 区分 `MXU0`..`MXU3`）。两个轴正交：bundle 槽选择*控制通道*，而 matmul opcode 的低两位 @ 89..90 选择*物理阵列*。重新实现者不能混淆“哪个 MXU slot”和“四个 MXU 中的哪一个”。`PufferfishTarget::MatrixStagingRegisterCount`（`0x1d4949e0`）返回 **1**，所以 PF 只有单个 MSR，因此没有 `Target` 字段，这与下面的 VF 不同。

---

## Viperfish (v5p) — 具名 Opcode 家族与 Latch 控制位

### MXU 控制区域

Viperfish 保留两个 `VectorExtended` 槽覆盖一个 operand pool 的模型以及 **−20** twin，扩宽到 64-byte bundle，并用*具名 opcode 家族*替换 Jellyfish 的 `GainLatchMode` enum。MXU 控制区域已从反编译的 `EncodeTensorCoreVectorExtended0PushmatrixBf16`（`0x1efaf820`）和 `MatrixMultiplyBf16` helpers 验证：

| 字段 | MXU0 abs | 宽度 | Op family | 值（Bf16） |
|---|---|---|---|---|
| MXU-id（unit） | 64 | 4 | always（先写入） | 0（MXU 0） |
| opcode-HIGH（matmul） | 57 | 7 | `MatrixMultiply<fmt>` | `0x1` |
| opcode-HIGH（latch/push） | 59 | 5 | `Pushmatrix<fmt>` | `0xe`（14） |
| data-format sub-disc | 51 | 4 | per-op | matmul Bf16 = 1 / push Bf16 = 3 |
| control（proto `+0x18`） | 48 | 3 | per-op | — |
| done-gains / latch flag | 55 | 2 | per-op | — |
| **Transpose** | 57 | 1 | `Pushmatrix*`（proto `+0x20`） | — |
| **Target** | 58 | 1 | `Pushmatrix*`（proto `+0x24`） | — |

latch（`Pushmatrix`）encode 从 `0x1efaf820` 逐字节读取为 `BitCopy(buf, 59, …, 5)`（opcode = `0xe`）、`BitCopy(buf, 51, …, 4)`（format = 3）、`BitCopy(buf, 48, …, 3)`（control）、`BitCopy(buf, 55, …, 2)`（done-gains/latch flag）、`BitCopy(buf, 57, …, 1)`（Transpose，proto `+0x20`）、`BitCopy(buf, 58, …, 1)`（Target，proto `+0x24`），再加上 abs 157/282/293/248/259/214/225/180（各 w6，proto 顺序）处八个共享 operand vregs。matmul helper `MatrixMultiplyBf16` @ `0x1efa2e40` 使用同一控制区域写入 `BitCopy(buf, 57, …, 7)`（opcode = `0x1`）、`BitCopy(buf, 51, …, 4)`（format = 1）、`BitCopy(buf, 48, …, 3)`（control）、`BitCopy(buf, 55, …, 2)`（done-gains），覆盖相同 operand pool。因此 matmul 和 push 共享每个字段位置，只在 opcode 宽度（7 vs 5）以及位 57/58 的重用上不同。所有偏移均为 **LSB-first**，**CONFIRMED**。

> **易错点 — opcode 字段在同一槽中会随 op 家族改变位置*和*宽度。** `MatrixMultiply<fmt>` 在 bit 57 写入 7-bit opcode（并在 bit 51 写入 4-bit `MatmulDataFormat`）；`Pushmatrix<fmt>` 在 bit 59 写入 5-bit opcode-HIGH（并在 bit 51 写入 4-bit Pushmatrix format）；`LoadMatrixRegister*` 复用 @ 57 的 7-bit 窗口，值为 `0x37`。在 push/latch 路径上，bits 57 和 58 被*重用*为 1-bit **Transpose** 和 **Target** 字段，也就是 matmul 路径上 7-bit opcode 的两个高位。decoder 通过 opcode-HIGH 值（push/latch `0xe` vs matmul `0x1`）区分两者，从而释放 LSB 区域用于 latch control。两个不同的 4-bit `data_format` enums 共享 abs 51：**matmul** `MatmulDataFormat`（Bf16 = 1，U8 = 2，S8 = 3，U4 = 4，S4 = 5，Bf8 = 6）和 **latch** Pushmatrix format（Rounded = 0，PackedIf8Conv = 2，Bf16 = 3，Bf8 = 4，U8 = 5，S8 = 6，U4 = 7，S4 = 8）。它们是同一字段中的不同 ordinal spaces。

### Transpose 和 Target Latch 字段

两个 1-bit latch 字段贯穿完整 producer → encoder → decoder → cost-model 链：

- **Transpose（abs 57）** = MLIR `tpu.matmul_push_rhs` 的 `transpose` 属性 — *在 RHS weight matrix 被 push 进 systolic array 时将其转置 latch*（latch 时完成 transposed-weights matmul）。该属性经 OpConversion 流向 `llo.vector_latch_i`，进入 `ViperfishTensorCoreEmitter::EmitVectorLatchCommon`（`0x141ba820`），其 Pushmatrix lambda 写入 proto `+0x20`，随后 `BitCopy(buf, 57, &proto[0x20], 0, 1)`。
- **Target（abs 58）** = `staging_register` 属性 = **MatrixStagingRegister (MSR)** bank select。`ViperfishTarget::MatrixStagingRegisterCount`（`0x1d49ace0`）返回 **2**，因此 1-bit 字段选择 MSR0 / MSR1。producer 将 `MatpushTarget` 参数写入 proto `+0x24`；`BitCopy(buf, 58, &proto[0x24], 0, 1)`。

`MatrixStagingRegisterCount` 在 Jellyfish（`0x1d490340`）和 Pufferfish（`0x1d4949e0`）上为 **1**，在 Viperfish（`0x1d49ace0`）和 Ghostlite（`0x1d497ae0`）上为 **2**；因此 JF/PF 不携带 `Target` 位（一个 MSR，无需选择），`Target` 字段仅从 v5p 起出现。解码后的 Transpose / Target 位会供给 cost model `SetReservations<MatpushModifier>`（`0x1c8abde0`，一个按 `{dtype × transpose}` 索引的 `array<int,19>`）以及 transpose-load latency `LatencyTableViperfish::XposeXLUReservationLatency`（`0x1c8a4f00`）。

> **易错点 — latch `transpose` 位*不是*独立 transpose op。** latch Transpose @ abs 57 在 weight matrix 被 push 进阵列时转置它（Pushmatrix-latch 属性）。另一个独立的 `EmitVectorTranspose<viperfish>`（`0x141c3f00`）op 家族会发出 `TransposeStart`/`Continue`/`End`/`Packed`/`Segmented` opcodes，通过 XLU/transpose unit 路由数据（PF opcode `0x40`/`0x48`）。它们是两种不同的 transpose 机制；重新实现者不能将其合并。

### Op-Family 名单（v5+）

每槽 94 个 `VectorExtended0` helpers 并不是 94 种不同布局；它们是一个 `{opcode, sub-format, MXU-id, operand}` 模板，按 opcode immediate 和 operand-present mask 特化：

| Op family | 变体 | 作用 |
|---|---|---|
| `MatrixMultiply<fmt>` | Bf16, Bf8, If8Bf16, S4, S8, U4, U8, F32Rounded | dense matmul step，每种 data format 一个 helper |
| `MatrixMultiply<fmt>Lgmr{Msra,Msrb}[Masked]` | 每 fmt × {Msra, Msrb} × {plain, Masked} | latch-via-LMR fused matmul（multi-pass K-tiling） |
| `Pushmatrix<fmt>` | Bf16, Bf8, S4, S8, U4, U8, Rounded, PackedIf8Conv | moving-operand push（matprep） |
| `LoadMatrixRegister{Gmr,Lmr}{Msra,Msrb}` | 4 | weight-stationary latch（opcode-HIGH `0x37`） |
| `*Transpose*`, `Segmented*`, `Packed*` | ~10 | systolic-array transpose |

`Masked` 和 `Lgmr{Msra,Msrb}` 后缀是 Jellyfish 6-entry `GainLatchMode → VEopcode` 表在 v5+ 上的实现形式：v3 上的小 enum 在 v5 上成为具名 opcode 家族。`Msra`/`Msrb` 选择 fused matmul 累加进入两个 MSR banks 中的哪一个（与 latch 路径上 `Target` 选择的 bank 相同）。

---

## GXC（v6e Ghostlite / v7 `6acc60406`）— 更宽 Opcode、FP8 与更大的 Twins

GXC 各代保留 v5+ codec 形状：两个 `VectorExtended` 槽覆盖一个 operand pool、`BitCopy` 打包、线性 `Opcode::Matches` 解码；但将 opcode 字段从 7→8 位加宽、移动控制区域，并重映射 dtype 集。

### Ghostlite（v6e / `glc`）

MXU opcode-HIGH 加宽为 **8 bits @ bit 58**（相对 Viperfish 的 7-bit @ 57），MXU-id 移到 bit 66，done-gains / latch flag 移到 bit 56，slot-encoder opcode bound 增长到 `0x70`（113 ops，相对 Viperfish 的 103）。dtype 集是完整 **8-format** `{F32, If8, Bf16, Bf8}`（float）+ `{U8, S8, U4, S4}`（int）。槽图见 [Ghostlite Bundle](bundle-gl.md#transformation-2--opcode-fields-widen-78-bits)。

latch / matmul opcode 是 @ abs 58 的*统一* 8-bit 字段，这在 encode 侧得到确认：`MatrixMultiplyBf16` @ `0x1f333ce0` 写入 `BitCopy(buf, 58, …, 8)`，值 `0x1`；`LoadMatrixRegisterGmrMsra` @ `0x1f33f140` 在*同一个* `BitCopy(buf, 58, …, 8)` 窗口写入 latch 值 `0x37`（55）。matmul opcode @ abs 58（w8）= `0x2`（Msra）/ `0x3`（Msrb），MSR-select = opcode LSB；decode-side `Opcode::Matches` 将低两位 58,59 == 3 分组为 `latch-class` 判别符，并在同一窗口读取 6-bit latch sub-opcode（14 = float / 15 = int），dtype-class @ abs 54（w2）选择 4-element class 内的子序号。`GhostliteTarget::MatrixStagingRegisterCount`（`0x1d497ae0`）= 2。MXU0↔MXU1 twin 是 **−21**：matmul opcode-HIGH 在 MXU0（VEx0）锚定于 abs 58，在 MXU1（VEx1）锚定于 abs 37 — `MatrixMultiplyBf16` VEx0 @ `0x1f333ce0`（opcode @ 58、fmt @ 52、control @ 49、done-gains @ 56）与 VEx1 @ `0x1f388440`（opcode @ 37、fmt @ 31、control @ 28、done-gains @ 35）相比，每个字段都精确偏移 21。此段所有偏移均为 **LSB-first**，且由 encoder `BitCopy` immediates **CONFIRMED**。

### `6acc60406`（v7 / `gfc`）

最新代是**仅浮点**：它移除了 integer matmul group，并支持四种 dtypes `{F32, E4m3, Bf16, E5m2}`，两个 FP8 formats 显式命名（相对 Ghostlite 的 `If8`/`Bf8`）。来自 `MatrixMultiplyBf16` VEx0 @ `0x1f99a920`（已验证 `BitCopy` immediates）：matmul opcode-HIGH 8-bit @ bit 62，data-format @ bit 57（w4），control @ bit 54（w3），done-gains / latch flag @ bit 61（w1）；MXU-id @ bit 70（w2），由 VEx0 dispatcher encoder `TensorCoreVectorExtended0Encoder::Encode` @ `0x1f996940` 写入（`BitCopy(buf, 70, …, 2)`）。八个 systolic source vregs 位于 bits 156 / 276 / 287 / 243 / 254 / 210 / 221（w6）和 47（w7）— 最后一个 operand 加宽到 7 bits。latch valid-guard 是对 bit 62 的单个 `bt`。见 [6acc60406 Bundle](bundle-gf.md#the-mxu-format-remap-and-fp8-gf-vs-ghostlite)。所有偏移均为 **LSB-first**，**CONFIRMED**。

MXU0↔MXU1 twin 是 **−25**，不是 −21：`6acc60406` 的 MXU0 控制区域比 Ghostlite 向高位漂移 +4 bits（matmul opcode-HIGH 58 → 62），而 MXU1 在两代中都锚定于相同的 abs 37，因此 MXU 间 delta 增大 4。encode-side `MatrixMultiplyBf16` VEx0 @ `0x1f99a920` 与 VEx1 @ `0x1f9d77e0` 显示每个字段都精确偏移 25：matmul opcode-HIGH 62 vs 37，format 57 vs 32，control 54 vs 29，done-gains 61 vs 36。

### 跨代字段摘要

MXU 控制区域，综合五代（matmul opcode-HIGH bit / latch opcode bit / inter-MXU twin）：

| 代 | Codename | Bundle | VE slots | Physical MXUs | Matmul opcode bit | Latch opcode bit | Twin | MSR count | dtype set |
|---|---|---|---|---|---|---|---|---|---|
| v2 | jellyfish | 41 B | 1 | 1 | （6-bit VEopcode @ 29..34，jump table） | — | n/a | 1 | `GainLatchMode 0..5` subset |
| v3 | dragonfish | 41 B | 1 | 2 | （相同 codec；mxu-id live） | — | n/a | 1 | 同 JF |
| v4 | pufferfish | 51 B | 2 | 4 | 91（9b，带 mxu-num @ 89..90） | 91（PushGains `0x20`..`0x34`） | −20 | 1 | 8（int + float） |
| v5p | viperfish | 64 B | 2 | 4 | 57（7b） | 59（Pushmatrix `0xe`） | −20 | 2 | 8（int + float） |
| v6e | ghostlite（`glc`） | 64 B | 2 | 2 | 58（8b） | 58（统一 8b，low-2 == 3） | −21 | 2 | 8（int + float） |
| v7 | `6acc60406`（`gfc`） | 64 B | 2 | 2 | 62（8b） | 62（统一 8b） | −25 | 2 | 4（仅 float） |

演进是一条清晰路径：单个 6-bit jump-table opcode（v2/v3）→ 携带物理 MXU 于低位的双槽 7/9-bit opcode（v4）→ 带独立 Transpose/Target latch 位的具名逐 dtype opcode 家族（v5p）→ latch 判别符折入 opcode 低位的更宽 8-bit 统一 opcodes（v6e）→ 仅浮点 FP8 重映射（v7）。systolic *contract* — weight-stationary array、matpush 8×128/4×256 tiles、共享 operand pool 上的 latch / push / matmul / matres 序列 — 始终不变；只有编码加宽、dtype 集转移，以及阵列维度从 128×128（JF–VF）跃迁到 256×256（Ghostlite、`6acc60406`）。

---

## Encode / Decode 路径

MXU 槽先构建为 proto `BundleSlot`，再序列化为字节：

```text
STAGE 1 — emit (build proto submessage)
  Emit{VectorMatmul,VectorLatch,VectorMatres}
    └─ CurrentBundle → GetPopulatedSlots (check slot free)
    └─ EmitVector{Extended,Result}Instruction  (create the proto submessage)
    └─ AddMxuNumTo*                              (stamp 2-bit MXU id, set present bit)
  result: a Bundle proto with slot_mask bit 0x080 (VE) / 0x100 (VR) set

STAGE 2 — encode (proto → raw bundle bytes)
  v3:   EncoderJf::EncodeBundleInternal (0x1e86c7c0)
          → zero buffer; splat kNeverExecute (31) into every slot predicate
          → slot_mask dispatch → EncodeVectorExtendedInstruction / EncodeVectorResultInstruction
  v4+:  TensorCoreCodecBase::Encode → per-slot <Slot>Encoder::Encode → BitCopy(buf, dst_bit, …)
  decode (inverse):
    v3   = DecoderJf::DecodeVectorExtendedSlot (0x1e854000): two-level jump table on the 6-bit opcode
    v4+  = TensorCoreVectorExtended{0,1}Decoder::Decode: staged byte-copy + linear Opcode::Matches sweep

两个 Jellyfish encoders 都会在写入任何槽前，将 kNeverExecute = 31（0xB834CFC）预填到每个槽的 predicate 字段中，因此 bundle 中缺失的 MXU op 会留下定义明确的 never-execute predicate，而不是垃圾。5-bit predicate 语义由所有槽共享：bits 0..3 选择 predicate register 0..14，值 15 = kAlwaysExecute（0xB834CF8），bit 4（值 16）是 predicate-negate flag，值 31 = kNeverExecute。kPredicateRegisterCount、kAlwaysExecute、kNeverExecute 直接从 ELF 的 0xB834CF4 读取为 15、15、31。

易错点 — empty 是 predicate-31，不是全零。 empty-slot mark 是非零戳记。重新实现若 memset bundle 为零并只填充 active slots，会使 inactive MXU slots 处于 predicate 0，即有效 predicate-register 引用，从而把空槽变成活跃的垃圾 matmuls。见 NOP / 未使用槽规范编码。

Systolic-Array 几何

数量	值	来源
Systolic array	128 × 128（JF/DF/PF/VF）/ 256 × 256（Ghostlite、`6acc60406`）	base `Target` `LaneCount`；`GhostliteTarget::MxuContractingSize`/`MxuNoncontractingSize` @ `0x1d497840`/`0x1d497860` = 256
MXUs per TensorCore	1（JF）/ 2（Dragonfish）/ 4（PF、VF）/ 2（Ghostlite、`6acc60406`）	`NumVexSlots`、`VectorIsa.mxu_count` f5
MXU-id 字段宽度	2 bits（JF/PF：physical select）/ 4 bits（VF MXU-id @ 64，glc @ 66）/ 2 bits（gfc @ 70）	encoder
matpush / latch tile	8 × 128 或 4 × 256	vmatpush perf-counter string
vreg shape	8 sublanes × 128 lanes	`Target::SublaneCount` @ `0x1d60f300` = 8，`LaneCount` @ `0x1d60f400` = 128

在 128×128 各代上，阵列由十六个 8×128 matpush tiles（或八个 4×256）填满；moving operand 流经阵列，一个 vmatmul 推进一个 systolic step，结果在阵列 latency 之后通过 vmatres 排出。per-format matmul latency tables 位于各代 cost 页面；它们不是跨代单一常量。

易错点 — 并非每代阵列都是 128×128。 Ghostlite 和 6acc60406 将 mxu_count 从 4 降到 2，但将 systolic dimension 翻倍到 256×256（C++ overrides GhostliteTarget::MxuContractingSize/MxuNoncontractingSize 返回 256；base Target 返回 128）。256 维度是 C++ 字面量，而不是 proto 字段；VectorIsa proto 只携带 lane_count = 128 和 mxu_count。为 v6e/v7 硬编码 128×128 阵列的重新实现会误算 latch tile 数量和每步吞吐量。见 Per-Codename HW Constants。

组件	关系
`EncoderJf::EncodeVectorExtendedInstruction` `0x1e869f00`	v3 MXU 槽字节 encoder（opcode @ 29..34）
`EncoderJf::EncodeVectorResultInstruction` `0x1e865ae0`	v3 matres 字节 encoder（predicate @ 22）
`JellyfishEmitter::EmitVectorMatmul` `0x140b92c0`	v3 matmul opcode dispatch（DoneWithGainsMode）
`JellyfishEmitter::EmitVectorLatch` `0x140b8c20`	v3 latch opcode dispatch（GainLatchMode table）
`TensorCoreCodecBase::Encode` `0x1d224300`	v4+ 通用 codec（BitCopy-packed slots）
`ViperfishTensorCoreEmitter::EmitVectorLatchCommon` `0x141ba820`	v5p Transpose/Target 字段 producer
`SetReservations<MatpushModifier>` `0x1c8abde0`	decoded transpose/target bits 的 cost-model consumer

交叉引用

Bundle Model — MXU 槽所在的 VLIW bundle、slot_mask dispatch 和 kNeverExecute 约定。
Jellyfish 41B Bundle — v3 VectorExtended/VectorResult 绝对位位置以及 12-byte-strip law。
Pufferfish 51B Bundle — v4 双 MXU −20 twin，以及携带物理 MXU number 的 9-bit matmul opcode。
Viperfish 64B Bundle — v5p MXU 控制区域、Transpose/Target latch 位和具名 opcode 家族。
Ghostlite Bundle — v6e 8-bit 统一 opcode 和 −21 twin。
6acc60406 Bundle — v7 仅浮点 FP8 重映射和 −25 twin。
Matprep / IAR / Latch — matpush WORD 表以及 latch/push operands 所经过的 IAR addressing。
Decode-Side: JF / PF 和 Decode-Side: VF / GXC — 按字节确认每个字段位置的 disassembler 逆过程。
InstBits Master DB 和 MC Emitter — 相同 vmatmul/vmatprep/vmatres MachineInstrs 的并行 LLVM-MC 编码路径。
LLO Opcode Enum — MXU mnemonics 所在的 LloOpcode 数值空间。
../cost/mxu-latency-overview.md — 消费 MXU-slot 字段的 cost model，以及各代 matmul latency tables。
../compiler/dot-conv-mxu-lowering.md — 填充 MXU 槽为 [matprep, latch, matmul…, matres] 序列的 HLO kDot / kConvolution descent。
Per-Codename HW Constants — 几何表引用的 mxu_count（1/2/4/4/2/2）以及 Ghostlite/6acc60406 的 256×256 systolic-dimension override。

MXU 槽 ​

摘要 ​

MXU Op 家族 ​

目的 ​

Op 名单 ​

VectorExtended 字段布局 ​

Latch Opcode 编码 ​

Systolic-Array 几何 ​

相关组件 ​

交叉引用 ​