Internal Wiki

主题

IarsPerTensorCore 与非 MXU 成本带 ​

地址适用于 libtpu-0.0.40-cp314 wheel 中的 libtpu.so（BuildID md5 89edbbe81c5b328a958fe628a9f2207d，未 strip，包含完整 C++ 符号）。其他版本会不同。下面每个整数都手工从 .rodata（LUT 和嵌入的 chip-parts proto）读出，并与 IDA 反编译交叉核对；状态见 Confidence 列。

摘要 ​

本页汇总三个按代际划分的成本模型量，它们共享同一个根：一个计数和一个分类器表面，并且都以 MXU 吞吐表使用的同一个 CycleTable::Instruction 枚举为键。第一，IarsPerTensorCore 值，即每个 TensorCore 的 Index-Add Registers（IAR0/IAR1）数量，它限制 matpush/indexed-load 寻址路径并输入成本模型。第二，Jellyfish/Dragonfish 上的非 MXU CycleTable 带（Instruction 0x11..0x20），即 vector-ALU、cross-lane（Xlu）和 EUP 成本单元；其取值在 MXU 延迟：JF / DF 固定，而其生产者在这里记录。第三，Viperfish matprep 阶段序号绑定，即把 matprep opcode 转换为 reservation 数组的机制。

统一事实是：IarsPerTensorCore 不是代码常量。它是 DWORD[Target+0x4a8]，由共享的 Target::Init 从嵌入的按代际 *_chip_parts.binarypb proto 的 VectorIsa 字段写入一次，并在运行时通过 embed:// 文件系统加载。从嵌入 blob 中提取后，各代的值都是 2；IAR 文件在代际间稳定（IAR0/IAR1），与 ISA 编码中 1-bit IarField 上限一致。因此上面的版本固定是必要但不充分的：如果某个芯片的 proto 不同，计数就会变化。

对于重新实现，契约是：

  • IarsPerTensorCore = DWORD[Target+0x4a8]，访问器 Target::IarsPerTensorCore (0x1d617280)；唯一写入者 Target::Init (0x1d60fc20) 来自 vector_isa() 字段 7；所有代的值都是 2。
  • JF/DF 非 MXU 带 Instruction 0x11..0x20 由 HLO 级成本模型（CostModel::RecordHloCycles + cost_model_util::Record*）作为以 HLO opcode / PrimitiveType / transfer role 为键的直接序号立即数发出；不存在用于它的 LLO→Instruction 分类器（MXU 分类器 CycleTableInstruction 只覆盖 0x00..0x10）。
  • VF matprep 阶段不携带独立分类器序号：matprep opcode 0x97..0x9a 在 GetViperfishInstruction 中 FATAL，reservation 由 MxuLatencyTable::GetResourceUsage 产生，并以 MatpushModifier{MatmulDataFormat, is_transpose, Msr} 为键。
IarsPerTensorCoreDWORD[Target+0x4a8] = 所有代 2（IAR0/IAR1，代际稳定）
访问器Target::IarsPerTensorCore 0x1d617280（return *(uint32*)(this+0x4a8)）
唯一写入者Target::Init 0x1d60fc20，来自 TpuSequencerParts::vector_isa() 字段 7
Proto 来源embed://tpu_chip_parts/<version>_chip_parts.binarypb, VectorIsa.f7
ISA 上限1-bit IarField (0x1ee3b380, >>13 & 1) ⇒ 2 个 IAR
非 MXU 带CycleTable::Instruction 0x11..0x20（vector-ALU / Xlu / EUP）
非 MXU 生产者CostModel::RecordHloCycles 0x130bbfe0 + cost_model_util::Record*（直接序号立即数）
VF matprep 绑定MxuLatencyTable::GetResourceUsage 0x1c8ae5c0, FlatHashMap<MatpushModifier, array<int,19>>
VF matprep opcode0x97..0x9a 在 GetViperfishInstruction 0x1c8a3300 中 FATAL

IarsPerTensorCore — 来源和按代际取值 ​

访问器和唯一写入者 ​

IarsPerTensorCore 是一个单指令访问器，读取 Target 对象偏移 0x4a8（298 * 4 = 0x4a8）：

c
// xla::jellyfish::Target::IarsPerTensorCore  @ 0x1d617280  (verified)
__int64 Target::IarsPerTensorCore(this) {
    return *((unsigned int*)this + 298);          // DWORD[Target + 0x4a8]
}
```text

`Target+0x4a8` 的唯一写入者是**共享的** `Target::Init`，没有任何按代际的 `*Target` 构造函数写入它。在 `Init` 的寄存器计数块中，该字段从 `VectorIsa` C++ 子结构（嵌入在 `TpuSequencerParts+0x1c`）读取，并受 VectorIsa-present 字节门控：

```c
// xla::jellyfish::Target::Init  @ 0x1d60fc20  (verified — the +0x4a8 store)
v93 = TpuSequencerParts::vector_isa(seqparts);    // = &seqparts[+0x1c]
if ( *(uint8_t*)(v93 + 0x18) == 1 ) {             // VectorIsa-present gate (seqparts+0x34)
    v94 = *(uint32_t*)(v93 + 0x14);               // IarsPerTensorCore  (seqparts+0x30 = VectorIsa.f7)
    *((uint32_t*)Target + 298) = v94;             // store at Target+0x4a8   ← the sole writer
    *((uint64_t*)Target + ...) = *(uint64_t*)(v93 + 0xc);   // adjacent MxusPerTensorCore qword
}

TpuSequencerParts::FromProto 将 proto VectorIsa 子消息（六个 int32 字段，proto 字段号 2..7）打包进 C++ 结构，其中字段 7 落在结构 +0x30 → Target+0x4a8。因此 IarsPerTensorCore 是 VectorIsa 字段 7 的 proto 值，由 TpuChipParts::DefaultsForVersion 通过 tsl::ReadBinaryProto("embed://tpu_chip_parts/<version>_chip_parts.binarypb") 在运行时加载。

提取出的按代际取值 ​

按代际的 blob 以未压缩形式嵌入 .rodata。从每个 blob 中解析 VectorIsa 子消息（f2=128, f3=8 前缀定位它；f7 是 IAR 计数）得到逐字节精确结果：

代际（代号）lane_count f2sublane_count f3mxu_count f5xlu_count f6iar_count f7 = IarsPerTensorCore
v2 Jellyfish1288112
v3 Dragonfish1288212
v4 Pufferfish1288422
v5p Viperfish1288432
v6e Ghostlite1288222
v7 6acc604061288222

所以每一代上 IarsPerTensorCore = 2，IAR 文件不会从 v2 到 v7 增长。这与 ISA 编码上限完全一致：SetIar slot 的 IarField 访问器（0x1ee3b380）是 >>13 & 1，单个 bit，只能寻址 IAR0/IAR1。VectorIsa 字段名不是推断出来的；TpuSequencerParts::FromProto（0x20b30700）按名称验证每个字段（"Invalid lane_count in vector_isa field"、"Invalid sublane_count …"、"Invalid mxu_count …"、"Invalid xlu_count …"、"Invalid iar_count …"），因此 f2=lane_count=128（所有代恒定）、f5=mxu_count（1/2/4/4/2/2）和 f6=xlu_count（1/1/2/3/2/2）是 CONFIRMED，而不只是位置推断。

注意 — 这个值是数据，不是代码。 访问器和 1-bit IarField 是不可变代码；取值是运行时 proto 字段。“所有代都是 2”的说法只对这个二进制中嵌入的 chip-parts blob 成立。重新实现必须从按代际 proto 读取它，而不是把它写死。IAR 值字段布局见 Chip-Parts Binarypb 和 ../isa/slot-matprep-iar-latch.md。

汇总的按代际寄存器文件表 ​

IarsPerTensorCore 是 Target+0x498..+0x4a8 处按代际寄存器文件块的最后一列，全部由同一个 Target::Init 从同一批 chip-parts blob 写入。四个寄存器计数从嵌入 proto 中逐字节精确重新推导：

GenSREG +0x498VREG +0x49cVMREG +0x4a0PREG +0x4a4IARS +0x4a8
v2 Jellyfish32328152
v3 Dragonfish32328152
v4 Pufferfish32328152
v5p Viperfish326416142
v6e Ghostlite326416142
v7 6acc60406326416142

（字段偏移逐字节精确来自 Target::Init + RegisterCount；取值来自嵌入的 chip-parts proto。SREG/VREG/VMREG/PREG 顺序遵循 Init 的 seq-0 type dispatch。）

JF/DF 非 MXU CycleTable 带（Instruction 0x11..0x20） ​

生产者 ​

MXU 分类器 CycleTableInstruction（0x1c89ca80）只映射 MXU 带：matmul opcode 0x8d..0x96 和 matprep/matpush opcode 0x9b..0xa5 到 CycleTable::Instruction 0x00..0x10。不存在用于序号 0x11..0x20 的 LLO→Instruction 分类器。这些序号由 HLO 级成本模型作为直接序号立即数（mov esi, 0xNN）发出，按 HLO opcode、PrimitiveType 或 memory-transfer role 为键，然后传给 GetCyclesForThroughput（JfCycleTable vtable +0x10）取得周期值，并传给 ResourceVector::Acc(Resource, cycles) 取得累加槽：

  • CostModel::RecordHloCycles（0x130bbfe0）— 每个 HLO 的入口；对 HloOpcode switch，并为逐元素 vector op 发出 0x12/0x13/0x14。
  • cost_model_util::Record* — RecordPackAndStCycles（0x13844740, 0x16）、RecordBroadcastSublaneChunkCycles（0x138448c0, 0x17/0x1b）、RecordMemXferCyclesImpl（0x13844e80）；Broadcast::RecordCostCycles（0x136e6940, 0x17）；Record{Conv,DepthwiseConv}KernelCycles / RecordConvolutionCycles / SpatialMajorConvolution::CalculateWindowMxuCycles 发射器（0x16/0x19/0x1c/0x1f）。

该带表 ​

offsetLUT（0xb438b70）和 resLUT（0xb438aec）已逐字节精确重读；周期值是在有效 mask 0x19FFC0821 下定价的 PerformanceJf 单元。整个带上 JF == DF（这些偏移都不是 DF override +0x28/+0x2c）。

InstrResResourceVector是否定价PerfOffJF/DF 周期发出上下文
0x11r6VectorEup否0x01未定价（SparseCore decomposer）
0x12r4VectorAlu1是0x33c1逐元素 vector-ALU（RecordHloCycles）
0x13r4VectorAlu1是0x3401逐元素 vector-ALU（RecordHloCycles）
0x14r3VectorAlu0是0x3441vector-ALU lane0（reduction-fn）
0x15r5VectorAluAny是0x39c1vector-ALU any（sublane reduce-window）
0x16r5VectorAluAny是0x3981pack-and-store / reduce-window / conv-kernel
0x17r2Xlu是0x9548broadcast / cross-lane（Broadcast::RecordCostCycles）
0x18r6VectorEup是0x3f81EUP / transcendental（RecordHloCycles）
0x19r5VectorAluAny是0x3681vector-ALU any（RecordConvolutionCycles）
0x1ar6VectorEup是0x3f41EUP / transcendental（RecordHloCycles）
0x1br2Xlu是0x9608cross-lane reduce-window
0x1cr2Xlu是0x94c8conv-kernel cross-lane
0x1dr2Xlu否0x01未定价
0x1er2Xlu否0x01未定价
0x1fr2Xlu是0x9588spatial-conv window-MXU
0x20r5VectorAluAny是0x39c1vector-ALU any（与 0x15 共享 0x39c）

四个 8-cycle Xlu 单元（0x17/0x1b/0x1c/0x1f）是 cross-lane 吞吐端口（broadcast-sublane、reduce-window、conv-kernel-window、spatial-conv-window）。1-cycle 单元是 vector-ALU（0x12..0x16、0x19、0x20）或 VectorEup（0x18/0x1a）。Res 列是 ResourceVector 槽 R[2]..R[6]（Xlu / VectorAlu0 / VectorAlu1 / VectorAluAny / VectorEup）。

怪癖 — Instr 0x15 和 0x20 共享偏移 0x39c。 两者读取同一个 VectorAluAny 单元（值为 1）；生产者选择发出哪个序号（sublane-reduce vs reshape/sparse-core），但它们定价相同。假设每个序号都有不同偏移的重新实现会分配一个冗余单元。

Viperfish Matprep 阶段序号绑定 ​

Matprep Opcode 在分类器中 FATAL ​

在 Viperfish 上，matprep opcode（0x97..0x9a）不会被分类为独立的 Instruction 序号，而是命中 GetViperfishInstruction 的 FATAL 分支（0x1c8a3300，opcode jump table 0xb43a104，默认分支 0x1c8a3e6a）。matmul opcode 0x9b 则读取 matmul_data_format() 并索引 VF matmul-format WORD 表 @0xa2d05c0，逐字节精确读为 {0xd4, 0xda, 0xf8, 0xfe, 0xe0, 0xe6, 0xec, 0xf2}（fmt 1..8 = f32/bf16/fp8e5m2/fp8e4m3/u8/s8/u4/s4）：

MatmulDataFormatdtypeVFinstr 序号（@0xa2d05c0）
1f320xd4
2bf160xda
3f8e5m2→bf160xf8
4f8e4m3→bf160xfe
5u80xe0
6s80xe6
7u40xec
8s40xf2

Reservation 表生产者 ​

matprep 阶段由 MxuLatencyTable::GetResourceUsage（0x1c8ae5c0）产生，VfCycleTable::GetCyclesForThroughput（0x1c89e2c0）对其进行包装。反编译确认：case 0 调用 viperfish::MxuLatencyTable::GetResourceUsage(&result, ..., 212, 3, 0)。GetResourceUsage 根据 Resource 参数和 matmul 序号分派，构造一个 Modifier 键，并在 FlatHashMap<Modifier, std::array<int,19>> 中查找；19-entry 数组由 MxuResource 索引（SetReservations<...> 写入 array[MxuResource]=val，边界 0x13 = 19）。matprep 4 阶段 systolic-feed 占用（r4..r7）是这 19 个 MxuResource 端口中的四个。

四种 modifier 键，每个都由 ViperfishPerformance 构造函数中的 SetReservations<...> 构造：

Modifier键字段作用
MatmulModifier{MatmulDataFormat}纯 matmul 步骤
MatpushModifier{MatmulDataFormat, is_transpose, Msr}matprep/matpush 阶段
MatresModifier(matmul-result key)matmul-result 阶段
VlxmrModifier(vector-load matrix-result key)vector-load matrix-result

MatpushModifier 就是 matprep-opcode → reservation 绑定：一个 matprep 阶段的 reservation 以 {MatmulDataFormat × transpose-of-gains × matpush-Msr-stage} → 19-entry MxuResource reservation 数组为键。matprep opcode 携带一个索引该表的 Modifier，而不是独立序号。这是 VF 对同一工作的实现；JF 将其折叠为扁平 offset-LUT 单元。见 Matmul-Mode Modifiers 和 MXU 延迟：VF。

怪癖 — matprep 成本表示跨代迁移。 JF 将 transpose-of-gains 折叠进 matmul opcode（扁平单元）；PF 将 matprep 折叠进 Latch op；VF 让 matprep opcode FATAL，并使用以 MatpushModifier 为键的 array<19> reservation；GL/GF 为每个 matprep 变体提供固定的 binary-search perf 行（带更小的 array<11> reservation）。跨代使用单一 matprep 成本模型会错误定价五代中的四代。见 ../isa/slot-matprep-iar-latch.md。

示例 ​

三个端到端 trace，分别覆盖上面的三个表面之一。

Viperfish 上的 IarsPerTensorCore ​

Viperfish 设备构造 ViperfishTarget，其基类链调用共享的 Target::Init（0x1d60fc20）。Init 已经取得 viperfish_chip_parts.binarypb TpuSequencerParts（通过 embed:// 加载）。它读取 vector_isa()+0x18（present 字节）即 1，然后读取 vector_isa()+0x14 = VectorIsa.f7 = 2，并将其存入 Target+0x4a8。后续 IarsPerTensorCore() 返回 2；IAR-setter helper 约束 iar_value < 2（IAR0/IAR1）。没有 ViperfishTarget 构造函数自己写入 +0x4a8；该值纯粹来自 proto 字段。

一个 JF/DF Broadcast 成本 ​

HLO broadcast 通过 CostModel::RecordHloCycles 路由到 Broadcast::RecordCostCycles（0x136e6940），后者发出序号立即数 Instr 0x17 并调用 GetCyclesForThroughput(0x17)。这读取 PerformanceJf[offsetLUT[0x17]] = PerformanceJf[0x954] = 8 cycles，然后调用 ResourceVector::Acc(r2 = Xlu, 8.0)。因此一个 JF/DF broadcast 成本为 8 个 cross-lane（Xlu）cycles。逐元素 add 则发出 Instr 0x12 → PerformanceJf[0x33c] = 1 cycle，落在 r4 = VectorAlu1。

一个 VF bf16 Matprep 阶段 ​

一个 bf16 kVectorMatmul（opcode 0x9b，MatmulDataFormat = 2）通过 GetViperfishInstruction 路由，其 0x9b 分支读取 WORD[0xa2d05c0 + 1*2] = 0xda（bf16 纯 matmul 序号）。它的 matprep 伴随项由 MxuLatencyTable::GetResourceUsage(0xda, Resource, gains)（0x1c8ae5c0）产生：0xda 分支为 matmul 构造 MatmulModifier{DataFormat = 2}，并为 prep 阶段构造 MatpushModifier{DataFormat, is_transpose, Msr = LatchOpcodeToMsr(0x97..0x9a)}，然后 FlatHashMap<MatpushModifier, array<int,19>>::find 返回 19-entry MxuResource reservation。matprep-stage 列 r4..r7 = {4, 12, 20, 28}，即四个 systolic-feed 阶段。bf16 matprep 阶段的成本是以 MatpushModifier 为键的 array<19>，绝不是独立序号。

相关组件 ​

ComponentRelationship
Target::IarsPerTensorCore 0x1d617280读取 DWORD[Target+0x4a8]（= 所有代 2）
Target::Init 0x1d60fc20从 VectorIsa.f7 唯一写入 +0x4a8
TpuChipParts::DefaultsForVersion 0x20b1b040加载 embed://...<version>_chip_parts.binarypb
CostModel::RecordHloCycles 0x130bbfe0发出非 MXU Instr 0x11..0x20 序号
CycleTableInstruction 0x1c89ca80仅 MXU 分类器（Instr 0x00..0x10）
GetViperfishInstruction 0x1c8a3300matprep opcode 0x97..0x9a 在这里 FATAL
MxuLatencyTable::GetResourceUsage 0x1c8ae5c0VF Modifier → array<19> matprep reservation

交叉引用 ​