ConvolutionCostState

本页中的所有地址都适用于 libtpu-0.0.40-cp314 wheel 中的 libtpu.so（BuildID md5 89edbbe81c5b328a958fe628a9f2207d）。该二进制未 strip，下面每个符号都是 demangled C++ 名称。Section map：.text/.rodata VMA == 文件偏移；.data.rel.ro VMA − 0x200000 == 文件偏移。所有地址都是 VMA。

摘要

ConvCostState 是成本模型在给 op 定价前，每个 convolution 构建一次的运行状态。它是 conv 形状上的 LLVM MachineFunction analysis cache 类比物：只提取一次 conv-like HLO，解码其 ConvolutionDimensionNumbers、kernel Shape 和 Window，把每个逻辑维度映射到其物理（layout-mapped）位置，在三个 operand shape 上运行 Target 的 per-axis chunk grid，并保存结果，供每个 kernel 定价叶子直接读取。该结构由 CostModel::GetConvolutionCostState @0x130a4b20 构建，并作为 const ConvCostState& 参数贯穿整个 conv-cost emitter 链。

实际上有两个不同的 struct，把它们混为一谈是核心陷阱。ConvCostState 是大的持久状态（1180+ bytes：+0x158/+0x1e8/+0x328 处的 protos、+0x350..+0x378 处的六个 layout-mapped dim size、一组 +0x380..+0x458 的 per-dim vector mask、+0x470..+0x484 的 21-byte ConvolutionLoweringStrategy dtype/pack flag，以及 +0x8/+0x78/+0xe8 的三个 inlined_vector<long,6> chunk-count grid）。ConvState 是小的 per-call dim-product struct（+0x0..+0x78），由 CostModel::RecordConvolutionCycles @0x130b6ce0 在写入过程中，在栈上根据 ConvCostState 的 chunk-count vector 除以 Target geometry 构建。定价叶子 CostModel::RecordConvKernelCycles @0x130caf20 从 ConvCostState 读取 dtype/pack flag，从 ConvState 读取 dim product，把它们乘成 op count，并把 cycles 写入 Matmul / Matpush / Xlu resource slot。

本页记录两个 struct 的字段偏移，以及把 conv 的 op count 转换成 cycles 的吞吐桥接。桥接是核心细节：conv 写入需要的 matmul-rate multiplier 是 VfCycleTable::GetCyclesForThroughput(CT 0)，它不是常量，而是一次对 MxuLatencyTable::GetResourceUsage(instr 212, res 3) 的调用，并且 res 参数是资源索引重映射（res 3 → array index 15 = MatmulAccA；res 11 → array index 0 = MatpushPushPort），不是 cycle seed。因此 VfCycleTable throughput integer 和 mxu-latency-overview 的 MxuLatencyTable reservation cycle 是从两个角度看到的同一个数字。

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

• ConvCostState 字段图：三个被复制的 proto、六个 logical→physical dim size、per-dim vector mask、ConvolutionLoweringStrategy flag bytes（只在 opcode 0x2b=convolution 时读取）以及三个 chunk-count vector。
• ConvState 字段图：+0x8..+0x78 的 dim product、+0x0..+0x6 的四个 shape-case flag byte，以及填充每个字段的 Product(chunk_vector)/SublaneCount/ChunksPerTile 公式。
• Matmul/Matpush/Xlu 写入：op count × thru(CT) × format-rate（0.5/0.25）÷ Target+0x4ac（matmul rate）或 Target+0x4b0（Xlu rate）÷ ElementPackingFactor。
• VfCycleTable → MxuLatencyTable 桥接及其 res→index 重映射。
• Target+0x4ac/+0x4b0 是 Target vector-ISA rate 字段，不是 ConvCostState 内部的 derate 字段。

Structs	xla::jellyfish::CostModel::ConvCostState（持久） · xla::jellyfish::CostModel::ConvState（per-call）
Builder	CostModel::GetConvolutionCostState @0x130a4b20 (3520 B)
ConvState 构建	CostModel::RecordConvolutionCycles (4-RV) @0x130b6ce0（构建于 rbp-0x1d0）
定价叶子	CostModel::RecordConvKernelCycles @0x130caf20
吞吐桥接	VfCycleTable::GetCyclesForThroughput @0x1c89e2c0 → MxuLatencyTable::GetResourceUsage @0x1c8ae5c0
thru(CT 0)	GetResourceUsage(instr 212, res 3, false) = array[15] = MatmulAccA = 8 (bf16)
thru(CT matpush)	GetResourceUsage(instr 267, res 11, false) = array[0] = MatpushPushPort = 2 (bf16)
Geometry divisors	Target::SublaneCount(), LaneCount(), ChunksPerTile()
Rate divisors	Target+0x4ac（matmul）、Target+0x4b0（Xlu）— vector-ISA fields

---

ConvCostState — 持久状态

用途

GetConvolutionCostState @0x130a4b20 的签名是 bool GetConvolutionCostState(const HloInstruction* inst, ConvCostState* out)。它提取 conv-like HLO（fusion_util::ExtractConvLikeHlo @0x1d6aa140；空结果会在 cost_model.cc:3374 的 conv_like_hlo != nullptr 处触发 LogMessageFatal），然后自上而下填充 struct：复制三个 shape/window proto，在三个 operand shape 上运行 ChunkCountsWithTmp，通过 LayoutUtil::MakeLogicalToPhysical 把每个逻辑 conv 维度映射到其物理位置，初始化 per-dim vector mask，把每个 dim（batch / outer / inner）分类进三个 mask vector，最后仅对真正的 convolution opcode 复制 ConvolutionLoweringStrategy flag block。

字段图

a3 是 ConvCostState*。偏移是从 struct base 开始的字节偏移；反编译器将它们渲染为 (char*)a3 + N（proto copies）或 *((_QWORD*)a3 + N/8)（scalar/pointer stores）。

Offset	内容	Store site / source
+0x0	const HloInstruction* conv	*(_QWORD*)a3 = ExtractConvLikeHlo(inst) @0x130a4b4f
+0x8	inlined_vector<long,6> — operand-shape chunk counts	ChunkCountsWithTmp(operand_shape) → Assign(a3+1) @0x130a4c25
+0x78	inlined_vector<long,6> — output-shape chunk counts	ChunkCountsWithTmp(output_shape) → Assign(a3+15) @0x130a4c38
+0xe8	inlined_vector<long,6> — kernel-shape chunk counts	ChunkCountsWithTmp(kernel_shape) → Assign(a3+29) @0x130a4c58
+0x158	ConvolutionDimensionNumbers	CopyFrom((char*)a3 + 344, …) @0x130a4b79
+0x1e8	Shape (operand/kernel shape)	Shape::operator=((char*)a3 + 488, …) @0x130a4b90
+0x328	Window	Window::CopyFrom((char*)a3 + 808, …) @0x130a4baa
+0x350	output-feature dim SIZE（physical）	logical→physical [+106] @0x130a51f7
+0x358	input-feature dim SIZE（physical）	logical→physical [+107] @0x130a5212
+0x360	output-spatial dim SIZE（physical）	logical→physical [+108] @0x130a5235
+0x368	input-spatial dim SIZE（physical）	logical→physical [+109] @0x130a5250
+0x370	kernel-feature dim SIZE（physical）	logical→physical [+110] @0x130a5279
+0x378	kernel-spatial dim SIZE（physical）	logical→physical [+111] @0x130a529b
+0x380	vector<bool> — per-dim mask A（batch-dim）	assign(N, 0) @0x130a52bb
+0x398	vector<bool> — per-dim mask B（outer-dim）	assign(N, 0) @0x130a52dd
+0x3b0	vector<bool> — per-dim mask C（inner-dim）	assign(N, 0) @0x130a52ff
+0x3c8 / +0x3e0 / +0x3f8	vector<long>（init = 1）	assign(N, 1) @0x130a531e/533d/535c
+0x410	vector<bool>	assign(N, 0) @0x130a537e
+0x428	vector<long>（init = 1）	assign(N, 1) @0x130a539d
+0x440 / +0x458	vector<long>（init = 0）	assign(N, 0) @0x130a53bc/53e2
+0x470..+0x484	ConvolutionLoweringStrategy (21 B)	GetConvolutionLoweringStrategy，仅当 opcode == 0x2b 时复制 @0x130a5705..577a

NOTE — 三个 chunk-count inlined_vector<long,6> 成员位于 +0x8/+0x78/+0xe8，stride 为 0x70（每成员 112 B，含 header）：三次 Storage<long,6>::Assign 调用在三次 ChunkCountsWithTmp 调用后立即分别指向 (__int64*)a3 + 1、+15、+29。第一个 vector 来自 operand shape（v128），第二个来自 output shape，第三个来自 kernel shape；如果忽略 operand vector，很容易绑定错。

ConvolutionLoweringStrategy Flag Block

+0x470..+0x484 区域是 mxu-latency-overview 兄弟页中的 21-byte ConvolutionLoweringStrategy。它会以一次 16-byte head 的 vmovups [r15+0x470] 加上 +0x47d 处的 qword tail 复制，仅当 conv-like HLO 的 opcode byte（[conv+12]）为 0x2b（convolution）时执行。RecordConvKernelCycles 逐字节读取它作为 dtype / packing 决策。struct 内部的字节偏移（由 @0x130caf20 的字节读取确认）：

CCS offset	Decimal	Flag（推断名称）	Read in RecordConvKernelCycles
+0x470	1136	perform_batch_packed_activation_reorg	== 1 gate @0x130cb040
+0x476	1142	perform_activation_if_bf16_packing	!= 0 / == 1 gate @0x130cb1b8/cb2bb
+0x477	1143	generate_combined_x8_packed_matmuls_and_latches	OR-combine @0x130caf38
+0x478	1144	generate_combined_x4_packed_matmuls_and_latches	OR-combine @0x130caf38
+0x479	1145	generate_post_xpose_x8_packed_matmuls	OR-combine @0x130caf38
+0x47a	1146	generate_bf16_packed_vmatmuls	^1 index-select @0x130cb37e
+0x47e	1150	perform_activation_bf16_packing	OR-combine @0x130caf38
+0x480	1152	generate_x8_packed_vlatches	== 1（0.25 quarter-rate gate）@0x130cb468
+0x481	1153	generate_x8_packed_vlatches_for_activations	OR-combine @0x130caf38
+0x482	1154	generate_x8_packed_vmatmuls	OR-combine @0x130caf38

这些 byte 会 OR-combine（@0x130caf38..af86）以决定是否查询 TransferSizeUtil::ElementPackingFactor @0x1d6b03e0（只有任一 flag 被置位时才调用），否则使用 rbx = 1 / 2 默认值。flag 名称来自 ConvolutionLoweringStrategy proto 字段顺序和 gate 语义推断；偏移和读取位置经过字节确认。

---

ConvState — Per-Call Dim-Product Struct

用途

ConvState 在 RecordConvolutionCycles (4-RV) @0x130b6ce0 内部构建于栈上（rbp-0x1d0），并以 const& 传给 RecordConvKernelCycles。它不会持久化。每个 dim 字段都是对 ConvCostState 的某个 chunk-count vector 做 Product(chunk_dim_sizes)，再除以一个 Target geometry 常量（SublaneCount() = 8、LaneCount() = 128、ChunksPerTile() = 16），并带有成对的 remainder / divisibility flag。三次 xla::Product 调用（@0x130b6f1d/6f43/6f69）分别归约 operand / kernel / output chunk vector；前四个 byte 是写入分支所依据的 conv-shape-case flag。

字段图

a3（在 RecordConvKernelCycles 中）是 ConvState*。反编译器把字段渲染为 a3[k]（8-byte stride）或 a3[byte]。

Offset	a3[k]	内容（推断角色）
+0x0	a3[0] byte	shape-case flag（"input == kernel" class）
+0x2	a3[2] byte	shape-case flag（"dim ≥ 2" class）
+0x3	a3[3] byte	feature-dim selector（作为 [+0x3]^1 index pick 使用）
+0x5	a3[5] byte	HasNonBatchDotBaseDilation result @0x130b79fc
+0x6	a3[6] byte	grouped-conv gate（== 1）
+0x8	a3[1]	spatial / output dim product
+0x10	a3[2]	spatial / output dim product（minor gate）
+0x18	a3[3]	spatial dim product
+0x20	a3[4]	output dim product
+0x40	a3[8]	feature-group / kernel-feature product（在 ShouldPackInputFeature 中读取）
+0x48	a3[9]	window-volume / spatial product
+0x50	a3[10]	the M dim
+0x60	a3[12]	feature dim product（与 +0x68 组成 [+0x3]^1-indexed pair）
+0x68	a3[13]	feature dim product（+0x60 的 sibling）
+0x70	a3[14]	output-spatial iteration count
+0x78	a3[15]	window-iteration count

(+0x60, +0x68) 这一对加上 [+0x3]^1 XOR-select（@0x130cafc3）是“哪个 feature dim 在前”的选择：写入会读取 a3[16*a3[3] + 104]（即 [+0x68] 区域）和 a3[16*(a3[3]^1) + 96]（即 [+0x60] 区域）。op count 是 a3[3] · a3[2] · a3[4] · a3[1]（[+0x18]·[+0x10]·[+0x20]·[+0x8] 块）× [+0x50] × feat0 × feat1，然后根据 flag bytes [+0x0]、[+0x2]、[+0x3] 分支折入 [+0x48] / [+0x70] / [+0x78]（@0x130cb011/0c0/0ee/113 的四个 conv-shape case），最后除以 ElementPackingFactor 和 ChunksPerTile。

NOTE — 每个 ConvState dim offset 的语义角色（哪个是 output-feature、input-feature、spatial、batch、kernel-window）来自 imul 链和 ChunkCountsWithTmp source order；这个 build 不带 DWARF，因此维度名称是解释，而偏移和公式形状直接来自反汇编。权威来源是 GetConvIterationCounts @0x130b3c80。

---

吞吐桥接 — VfCycleTable → MxuLatencyTable

桥接为何存在

RecordConvKernelCycles 需要 matmul 和 matpush 的吞吐（cycles-per-op），而不是基础 Performance grid 的 per-op latency。这个吞吐就是 mxu-latency-overview 中的 per-resource occupancy。VfCycleTable::GetCyclesForThroughput(CycleTable::Instruction) @0x1c89e2c0 是适配器：它把每个 32-entry CycleTable::Instruction（CT）ordinal 映射到 MXU occupancy table 中的正确调用。conv 写入通过 Target 的 CycleTable vtable 到达它，也就是 @0x130caf20 中的 (*(…+16LL))(this, CT, …) 间接调用（例如 @0x130cb... 的 CT 0、@0x130cb375 的 CT 1）。

CT → (instr, res, transpose) 表

@0x1c89e2c0 的反编译（对 CT 执行 switch(a2)，default → return 1）确认了三种调用形态：matmul/matpush CT 调用 MxuLatencyTable::GetResourceUsage；mxres/Xlu CT 调用 ViperfishPerformance::GetResourceUsage(…, res 14) + 1；CT 10/16 在 cycle_table.cc:682 处执行 LogMessageFatal("Unsupported PushGainsS4.")。

CT	Call	角色
0	MxuLat.GetResourceUsage(212 matmul', res 3, false)	conv matmul-rate thru(CT 0)
1	MxuLat.GetResourceUsage(218 matmul'', res 3, false)	matmul-rate（alt）
4	MxuLat.GetResourceUsage(230 matmul, res 3, false)	matmul-rate（base）
5	MxuLat.GetResourceUsage(267 matpush, res 11, false)	matpush-rate
6	MxuLat.GetResourceUsage(271 matpush', res 11, false)	matpush'-rate
9	MxuLat.GetResourceUsage(277 matpush'', res 11, false)	matpush''-rate
10, 16	LogMessageFatal("Unsupported PushGainsS4.")	abort
11	MxuLat.GetResourceUsage(267, res 11, **true**)	matpush-rate（transposed）
12	MxuLat.GetResourceUsage(271, res 11, **true**)	matpush'-rate（transposed）
15	MxuLat.GetResourceUsage(277, res 11, **true**)	matpush''-rate（transposed）
23	VfPerf.GetResourceUsage(296, res 14) + 1	mxres / Xlu（matres TC）
27	VfPerf.GetResourceUsage(302, res 14) + 1	mxres
28	VfPerf.GetResourceUsage(287, res 14) + 1	conv MXRES（Xlu）
29	VfPerf.GetResourceUsage(291, res 14) + 1	mxres
31	VfPerf.GetResourceUsage(298, res 14) + 1	mxres
all others	return 1	vector / matprep-bucket default

transpose flag 是第 4 个 GetResourceUsage 参数；matpush key 的预变换（267 direct、271 → latch_mode ^ 0xB、277 → latch_mode | 0x14）在 GetResourceUsage 内部应用（见 mxu-latency-overview）。

吞吐桥接 — res 是 Resource-Index Remap

MxuLatencyTable::GetResourceUsage(instr, res, transpose) @0x1c8ae5c0 不会直接用 res 索引其 array<int,19> reservation vector。它会先重映射：

function GetResourceUsage(instr, res, transpose):       // VF @0x1c8ae5c0
    if   res == 3:  idx = 15        // → MatmulAccA
    elif res == 11: idx = 0         //  → MatpushPushPort
    else: return InvalidArgument("Unsupported kind of resource")   // mxu_latency_table_vf.cc:547

    // build the modifier key from the opcode, then find:
    switch (instr):
        case 212: key = MatmulModifier{format=1}; map = this+0x20
        case 218: key = MatmulModifier{format=2}; map = this+0x20
        case 230: key = MatmulModifier{format=6}; map = this+0x20
        case 267: key = MatpushKey(latch_mode);          map = this+0x00
        case 271: key = MatpushKey(latch_mode ^ 0xB);    map = this+0x00
        case 277: key = MatpushKey(latch_mode | 0x14);   map = this+0x00
        default:  LogFatal("Unsupported opcode")          // vf.cc:578

    entry = map.find(key)                 // miss → out_of_range
    array<int,19> v = entry.value         // vmovups: matmul reads [rdx+8], matpush reads [rdx+4]
    CHECK(idx < 0x13)                      // == 19
    return (float) v[idx]                  // the hold-cycle count
```text

因此 `res` 参数 3/11 是**命名子资源 selector**，会重映射为具体 array index（`3 → 15 = MatmulAccA`，`11 → 0 = MatpushPushPort`），然后*那个*索引才访问 per-modifier reservation array。因此 matmul CT（0、1、4）读取 `array[15]`；matpush CT（5、6、9、11、12、15）读取 `array[0]`。`VfCycleTable` throughput integer **就是** `MxuLatencyTable` reservation cycle，这两个成本表使用的是同一组数字。

| `thru(CT 0)` — matmul rate（`array[15]` = `MatmulAccA`，按 `MatmulModifier` format） | value |
|---|---:|
| bf16（format group 1） | 8 |
| format group 2 | 16 |
| int8 / wide（format group 4） | 32 |

| `thru(CT 5)` — matpush rate（`array[0]` = `MatpushPushPort`，按 `MatpushModifier`） | value |
|---|---:|
| bf16 narrow | 2 |
| bf16 mid | 4 |
| int8 / x8 wide | 8 |

> **NOTE —** 非 bf16 条目遵循 per-format reservation group；每个 `MatmulDataFormat` 对应的 live `MatmulModifier`/`MatpushModifier` map 给出精确值。`{2,1,1}` / `{4,3,2}` / `{8,7,6}` reservation triplet 的**第一个**元素（`MatpushPushPort` value）正是 `thru(CT matpush)`；另外两个（matrix-staging-register holds）由 MXU scheduler 中单独的 `GetResourceUsage` 调用读取，而不是由这个桥接读取。

### 每代桥接

该桥接按 generation 区分。Ghostlite helper `GlcCycleTable::GetCyclesForThroughputHelper` `@0x1c89ed20` 结构相同，但使用 Ghostlite instruction set：matmul CT（0、1、4）使用 instrs `292`/`298`/`310`（`0x124`/`0x12a`/`0x136`）；matpush CT（5、6、9、10）使用 `347`/`349`/`356`/`358`（`0x15b`/`0x15d`/`0x164`/`0x166`）；GL `MxuLatencyTable` `array<int,11>` 重映射 `res 4 → idx 3` 和 `res 9 → idx 9`。`GfcCycleTable::GetCyclesForThroughputHelper` `@0x1c89f400` 是 TPU7x 变体，形状相同（`array<int,11>`）。重新实现必须从 per-gen helper 读取 (instr, res) pair。

---

## 写入 — 组合起来

`RecordConvKernelCycles` `@0x130caf20` 从 `ConvState` 计算一个 op count，并把 cycles 写入三个 slot（`ResourceVector::Acc(slot, cycles)`）：

| Slot | `Acc` index | 写入 | Site |
|---|---:|---|---|
| `Matmul` | `R[1]` | `op_count × thru(CT 1) × 0.5 ÷ Target+0x4ac ÷ ElementPackingFactor` | `@0x130cb...`（Acc `@line 560`） |
| `Matpush` | `R[0]` | `matpush_count × thru(CT matpush)` | Acc `@0x130cb...`（`@line 609`） |
| `Xlu` | `R[2]` | `thru(CT mxres) × ChunksPerTile × remainder ÷ Target+0x4b0` | Acc `@line 259/428` |

当 format flag `v166` 为 clear 时，matmul 会乘以半速 `0.5`（`.rodata 0xa2df5c8`）；只有当 `ConvCostState+0x480`（`generate_x8_packed_vlatches`）`== 1` **且** `Target+0x398` `== 2`（`@0x130cb468/cb471` 的 topology gate）时，才会替换为 `0.25`（`.rodata 0xa2df6d0`）。matmul-rate divisor 读取为 `vcvtsi2sd xmm1, [rax+0x4ac]`（`Target+0x4ac`）；Xlu divisor 则通过 `CycleTable` 的 `Target` 指针到达 `[rax+0x4b0]`。

> **GOTCHA —** `Target+0x4ac`（matmul rate）和 `Target+0x4b0`（Xlu rate）是 **`Target` 字段**，不是 `ConvCostState` 字段。它们是 chip-wide vector-ISA throughput-rate 参数，在 `Target::Init` 中从 `TpuSequencerParts::vector_isa()` `@0x20b31840` 填充：`vector_isa[+0xc]`（8 bytes）→ `Target+0x4ac` & `+0x4b0`；`vector_isa[+0x14]` → `Target+0x4a8`。conv 写入通过 `*(_QWORD*)a1`（`CostModel` 的 `Target*`）以及 `CycleTable` 的 `Target*` 读取它们，从不从 `ConvCostState` 读取；它们不是 conv-local derate。

由于 `Matmul`、`Matpush` 和 `Xlu`（`R[0]`/`R[1]`/`R[2]`）位于 `MaxResourceCycles` 的普通-MAX 组中（MXU pipes 会重叠，见 [`resource-enum`](resource-enum.md)），compute-bound conv 的 bundle 成本是 `≈ max(Matmul, Matpush, Xlu)`，不是它们的和。

---

## 示例 — Viperfish 上的 bf16 Conv

```text
thru(CT 0)  = VfCycleTable::GetCyclesForThroughput(0)
            = MxuLatencyTable::GetResourceUsage(instr=212, res=3, false)
            = MatmulModifier{format=1 bf16}, array[remap(res=3)=15]
            = array[15] = MatmulAccA = 8

R[1] Matmul  = op_count × 8 × 0.5 ÷ Target[+0x4ac] ÷ ElementPackingFactor[bf16=1]

thru(CT 5)   = GetResourceUsage(instr=267, res=11, false) = array[0] = MatpushPushPort = 2
R[0] Matpush = matpush_tile_count × 2

thru(CT 28)  = VfPerf.GetResourceUsage(287, res=14) + 1            (the Xlu mxres-read rate)
R[2] Xlu     = thru(CT 28) × ChunksPerTile × remainder ÷ Target[+0x4b0]

bundle cost ≈ max(R[0], R[1], R[2])                                (MXU pipes overlap)

对 int8-x8 conv，matmul format 选择 group 4 → array[15] = 32（matmul）、array[0] = 8（matpush），即 bf16 rate 的 4×，对应 four-byte-plane x8 latch sequence。

---

兄弟项 — Reduce-Window Cost 复用 ConvState

CostModel::RecordReduceWindowCycles @0x130b5ec0 是 pooling 兄弟项。它直接接收三个 Window（不是 ConvCostState），并通过同样的 Product / SublaneCount / ChunksPerTile division 在 rbp-0x1d0 构建相同的栈上 ConvState。一个简单路径是：当 CostModel+0x14 为 false 时，它向 R[7] VectorLoad、R[3] VectorAlu0、R[4] VectorAlu1、R[2] Xlu 写入零并返回（@line 782..788 的 Acc(a9, 7/3/4/2)）。否则它按 GetReduceWindowType @0x1454d4a0 分派（0 → RecordLaneReduceWindowCycles，1 → RecordSublaneReduceWindowCycles，2 → RecordMajorReduceWindowCycles），Lane 和 Sublane 路径受 !HasBaseDilation gate 控制。lane 叶子（@0x130c97e0）写入 R[7] VectorLoad（op_volume = ConvState[+0x20]·[+0x18]·[+0x70]）、每元素 combine 的 R[5] VectorAluAny（thru(CT 22)）、UpdateCostBasedOnReductionFunction combiner cost（count = op_volume × (ConvState[+0x68] − 1)），以及跨 lane drain 的 R[2] Xlu（thru(CT 27) ÷ Target+0x4b0）。完整 pooling cost 在 reduce-window-pooling-cost。

---

函数图

Function	Address	角色
CostModel::GetConvolutionCostState	0x130a4b20	构建 ConvCostState (3520 B)
CostModel::RecordConvolutionCycles (4-RV)	0x130b6ce0	构建 ConvState，调用 kernel leaf
CostModel::RecordConvKernelCycles	0x130caf20	dim → MXU 写入（R[0]/R[1]/R[2]）
VfCycleTable::GetCyclesForThroughput	0x1c89e2c0	CT → MXU-occupancy adapter
viperfish::MxuLatencyTable::GetResourceUsage	0x1c8ae5c0	res→index remap + array[idx]
GlcCycleTable::GetCyclesForThroughputHelper	0x1c89ed20	Ghostlite bridge（array<int,11>）
GfcCycleTable::GetCyclesForThroughputHelper	0x1c89f400	TPU7x bridge（array<int,11>）
fusion_util::ExtractConvLikeHlo	0x1d6aa140	conv-like HLO extractor
convolution_util::GetConvolutionLoweringStrategy	0x13192820	21-byte flag block
Target::ChunkCountsWithTmp	0x10c8b6e0	per-axis chunk grid → +0x8/+0x78/+0xe8
TransferSizeUtil::ElementPackingFactor	0x1d6b03e0	per-dtype packing divisor
ResourceVector::Acc	0x1c89adc0	slot 写入

---

相关组件

Name	关系
mxu-latency-overview	桥接读取的 MxuLatencyTable 模型，即 thru(CT) 返回的 reservation array
vf-cycletable	32-entry CycleTable::Instruction dump，以及每代完整 throughput bridge
window-description-cost	向 ChunkCountsWithTmp 提供输入的 conv/DMA byte+throughput primitive
reduce-window-pooling-cost	复用相同 ConvState build 的 pooling 兄弟项
resource-enum	写入所访问的 23-slot ResourceVector；MaxResourceCycles overlap model

---

交叉引用

• MXU Latency Overview — MxuLatencyTable reservation model、MatmulModifier/MatpushModifier key，以及此桥接暴露的 array[idx] 读取
• VfCycleTable — 32-entry CT → (instr, res) dump，以及完整的每代 GetCyclesForThroughput bridge
• WindowDescription Byte-Cost — conv/DMA byte + throughput primitive
• Reduce-Window / Pooling Cost — RecordReduceWindowCycles，复用 ConvState 的 pooling 兄弟项
• Resource Enum (23-slot) — conv 写入访问的 ResourceVector slot（Matmul/Matpush/Xlu）以及 MaxResourceCycles overlap rule
• MXU Slot — conv 写入定价的 LLO MXU instructions（matmul / matpush / matres）