感知 Bundle 的成本
本页中的每个偏移、数值和地址都按字节精确地读取自
libtpu-0.0.40-cp314wheel 中的libtpu.so（BuildID md589edbbe81c5b328a958fe628a9f2207d）。其他版本会不同。所有.rodata地址都是虚拟地址；对这个二进制而言，.rodataVMA == 文件偏移（section[11]at0x84a0000），且.textVMA == 文件偏移。
摘要
TPU TensorCore 每个周期发射一个 VLIW bundle，把每个已填充的 slot 同时发往不同的功能单元。成本模型也按此给 bundle 定价：它不会把各 slot 的 op 成本相加，而是取它们的最大值，也就是最慢的已占用功能单元 lane 会阻塞发射，其他 lane 都免费重叠。本页记录这个最大值如何计算（ResourceVector::MaxResourceCycles），每个 op 的吞吐周期如何在归约前存入各 lane 累加器（通过 AccumulateInstructionUsage 调用 Acc），当多个 op 或子计算被打包在一起时两个成本向量如何组合（Add），每次迭代成本如何按循环 trip count 缩放并摊销 DMA 启动成本（ScaleResource / GetSubset），以及最终的 bundle 周期数如何进入依赖延迟轴（LatencyBetween）和 LatencyHidingScheduler。
这是 slot 模型在 bundle 层级的消费者。23-slot ResourceVector、Acc 写入，以及 MaxResourceCycles 归约的三组结构在 Resource Enum 页面介绍；被写入的各类吞吐整数来自 CycleTable / Per-Opcode Cycle Constants 页面。本页把它们连接起来：展示完整的 op 级 → bundle 级 → loop 级 → schedule 级流水线，在每一步用字节地址锚定，并固定归约所编码的重叠语义。
熟悉 LLVM 的读者可以保留一个类比和一个差异。类比是：MaxResourceCycles 相当于 TPU 版的 LLVM MCSchedModel 资源压力查询，问题是“给定这个 group 消耗的 ProcResource 占用，瓶颈周期数是多少？”差异是：这里没有 scoreboard，也没有乱序窗口，bundle 字本身就是发射包，因此每个 bundle 只在固定的 23-slot 向量上取一次 max；唯一会串行而非重叠的资源是内存传输项和双 lane vector-ALU 端口平衡。其他所有内容，包括连续的 MXU push，都重叠。
对重新实现而言，契约是：
- Bundle 成本是在各 lane 占用上取
max，不是sum。MaxResourceCycles用三条规则把 23-slotResourceVector归约为一个标量：vector-ALU 三元组上做 50% 端口平衡混合，四个内存传输项做串行求和，其他所有项（MXU、XLU、vector ports、ICI、SparseCore）做普通max。 - 每个 op 的写入先于归约。 每个 LLO op 把
(double)GetCyclesForThroughput(class)加入 slotGetResource(class)；bundle 的向量是其中各 op 写入的按 slot 求和，之后MaxResourceCycles再将其折叠。 - 跨向量组合会 ADD 每个 slot，除了两个 DMA-latency slot 用
max组合。Add（以及 loop-prologue 的GetSubset）把 DMA 启动视为只付一次，把 DMA 带宽和计算视为按单位付费。 - 循环缩放按 trip count 乘以每个 slot 成本，但跳过两个 DMA-latency slot（
ScaleResource），因此软件流水化循环只支付一次 DMA 启动，而每次迭代支付带宽/计算。 - 延迟轴是独立的，也按每个 op-pair 用
max组合。LatencyBetween返回两个LloValue之间的读后写停顿，它是在适用的每代 latency-array 字段上取 max，并带硬性下限；调度器组合吞吐成本（本页）和依赖延迟来把 op 放进 bundle。
| Bundle 归约 | ResourceVector::MaxResourceCycles() @ 0x1c89b9e0 — 23-slot → 标量 max |
| 每 op 写入 | ResourceVector::Acc(Resource, double) @ 0x1c89adc0; 驱动 AccumulateInstructionUsage::operator() @ 0x144fd720 |
| 跨向量组合 | ResourceVector::Add(rv, AddOptions) @ 0x1c89b820; 默认值 AddOptions::Defaults @ 0x1c89ada0 |
| 循环缩放 | ResourceVector::ScaleResource(Resource, double, ScaleOptions) @ 0x1c89b6a0; 默认值 ScaleOptions::Defaults @ 0x1c89ad80 |
| 循环子集 | ResourceVector::GetSubset(SubsetOptions) @ 0x1c89bb00 |
| Bundle 组装 | CostModel::GetCycles @ 0x130aade0 — GetHloResourcesImpl + Add |
| 延迟轴 | LatencyTable::LatencyBetween(LloValue,LloValue) @ 0x1c89f820; JF LatencyBetweenInternal @ 0x1c8a0d60 |
| 调度入口 | CostModelLatencyEstimator::GetLatencyBetween(HloGraphNode,HloGraphNode) @ 0x10ff8f00 |
| 混合常量 | qword_A2DF5C8 = 0.5（vector-ALU 端口平衡因子） |
三层成本流水线
一个已调度区域的成本自底向上分三层计算，每层都有自己的组合规则。重新实现者必须区分它们：在二进制使用 max 的地方使用 sum，或者缩放二进制保持不变的 slot，都会给整个 op 类定错价。
op level ── per LLO op: Acc(GetResource(class), GetCyclesForThroughput(class))
│ into one shared 23-slot ResourceVector (slot-wise SUM across ops)
▼
bundle level ── MaxResourceCycles(rv) → ONE scalar = slowest occupied lane
│ (plain MAX, except VectorAlu blend + MemXfer serial sum)
▼
loop level ── Add(rv_iter) combine two regions (per-slot SUM, DMA-latency MAX)
ScaleResource(rv, trip) scale per-iteration (per-slot ×trip, DMA-latency skipped)
GetSubset(rv, opts) prologue/tail split (keep input-DMA / keep compute+output)
│
▼
schedule ── LatencyBetween(a,b) RAW dependency stall (per-op-pair MAX over latency fields)
+ bundle cycle cost → LatencyHidingScheduler list priority
```text
> **NOTE —** op 级写入和 bundle 级归约使用*同一个* `ResourceVector` 对象：op 累加到 slot 中（`+=`），然后整个向量只归约一次。向量绝不会按 op 单独归约。因此，不同单元上的两个 op 成本是两者周期的 max（一个 bundle），而同一单元上的两个 op 成本是它们的和（该 slot 累加了两者）。slot 布局见 [Resource Enum](resource-enum.md)。
---
## 每 op 写入 — 通过 `AccumulateInstructionUsage` 调用 `Acc`
在任何归约之前，每个 LLO op 的吞吐周期都会落入恰好一个 slot。规范驱动是 SDC sequence-checker 用来重新推导 bundle usage 的 lambda；它是最清晰的、gen-invariant 写入路径反编译，字节精确：
```c
// AccumulateInstructionUsage(...)::$_0::operator() @ 0x144fd720 (decompiled, exact)
int operator()(CycleTable *ct, ResourceVector *rv, CycleTable::Instruction cls) {
Resource r = CycleTable::GetResource(ct, cls); // op→slot (flat LUT 0xb438aec)
int64 cyc = ct->vtable[+0x10](ct, cls); // GetCyclesForThroughput(cls)
ResourceVector::Acc(rv, r, (double)cyc); // rv[r*8] += cyc
return 1;
}这里固定了三点。(1) op→slot 映射 GetResource 是单个平坦 .rodata 查表，在各代之间相同（CycleTable Family）。(2) 吞吐周期是 vtable slot +0x10 上的每代 virtual（GetCyclesForThroughput）。(3) 写入是 Acc，即对命名 slot 做 +=，不是覆盖，因此同一 lane 上的多个 op 会累加。Acc（0x1c89adc0）会边界检查 resource < 0x17（23 slots），失败时用 trapping ud1，并按 8-byte double stride 执行 rv[resource*8] += cycles。
对非 MXU op，emitter 的 Record* 路径会直接写入：cost_model_util::RecordInputMemXferCycles / RecordOutputMemXferCycles 写入四个 MemXfer* slot（R[9..12]），RecordHloCycles（0x130bbfe0）为 HLO 级定价包装 Acc(GetResource, GetCyclesForThroughput) 路径，matmul/matprep 家族则经过每代 MxuLatencyTable。它们全部写入同一个稍后由 MaxResourceCycles 归约的 ResourceVector。
Bundle 归约 — MaxResourceCycles
MaxResourceCycles @ 0x1c89b9e0 是感知 bundle 成本的核心。它接收填好的 23-slot ResourceVector，返回一个 double：bundle 的发射成本（周期）。反编译是一个平坦汇编块；缩减后是三条规则：
// xla::jellyfish::ResourceVector::MaxResourceCycles() @ 0x1c89b9e0 (decompiled, exact shape)
double MaxResourceCycles(ResourceVector *rv) {
double a = rv[3], b = rv[4], c = rv[5]; // VectorAlu0 / VectorAlu1 / VectorAluAny
if (c > 0.0) { // (A) 2-lane vector-ALU port balance
double d = min(a - b, c); // move "any" work to the less-busy lane
c -= d; b += d; // (vsubsd/vminsd/vsubsd/vaddsd, [+0x18..+0x28])
c *= 0.5; // residual overlaps at 50% (qword_A2DF5C8)
a += c; b += c;
}
double acc = max(a, b); // vec_alu bottleneck
double mem = rv[9] + rv[10] + rv[11] + rv[12]; // (B) MemXfer serial sum [+0x48..+0x60]
acc = max(acc, mem);
// (C) plain MAX over every other slot, in physical order:
acc = max(acc, rv[0]); // R[0] Matpush [+0x00]
acc = max(acc, rv[1]); // R[1] Matmul [+0x08]
acc = max(acc, rv[2]); // R[2] Xlu [+0x10]
acc = max(acc, rv[6]); // R[6] VectorEup [+0x30]
acc = max(acc, rv[7]); // R[7] VectorLoad [+0x38]
acc = max(acc, rv[8]); // R[8] VectorStore [+0x40]
for (r in 13..22) // R[13..18] ICI, R[19..21] SparseCore, R[22] reserved
acc = max(acc, rv[r]); // [+0x68..+0xb0]
return acc;
}
```text
反编译会按这个物理顺序读取 slot：`[+0x18]`/`[+0x20]`/`[+0x28]` 是混合三元组，`[+0x48]`+`[+0x50]`+`[+0x58]`+`[+0x60]` 是内存求和，然后是一串对 `[+0x00]`、`[+0x08]`、`[+0x10]`、`[+0x30]`、`[+0x38]`、`[+0x40]` 以及 `[+0x68]` 到 `[+0xb0]` 的 `vmaxsd`。`0.5` 是字面量 `qword_A2DF5C8`。
这三条规则编码了三种硬件重叠模型：
| 组 | Slots | 规则 | 硬件含义 |
|---|---|---|---|
| (A) vector-ALU | `R[3]`,`R[4]`,`R[5]` | 端口平衡 + 50% residual | 两条专用 ALU lane + 一条会负载均衡的 "any" lane；无法平衡的 residual 以一半重叠 |
| (B) memory | `R[9]`,`R[10]`,`R[11]`,`R[12]` | 串行 **sum** | input-DMA startup + input bytes + output-DMA startup + output bytes 顺序发生 |
| (C) everything else | `R[0..2]`,`R[6..8]`,`R[13..22]` | 普通 **max** | 独立功能单元（MXU pipes、XLU、vector ports、ICI links、SparseCore）完全重叠 |
> **GOTCHA — MXU pipes 位于普通 `max` 组。** `R[0]`（Matpush）和 `R[1]`（Matmul）会重叠，因此同一个 bundle 中连续的 MXU gain-push 和 matmul issue 成本是其周期的 max，而不是和。把 MXU op 串行化的成本模型会显著高估每个 matmul-bound bundle；在 `6acc60406` 上，单个 bf16 matmul/latch 成本为 212 cycles（[Per-Opcode Cycle Constants](per-opcode-cycle-constants.md)），相加得到 424，而硬件（以及此模型）收取 212。相信 max。
>
> **GOTCHA — 只有内存传输会串行化。** 整个归约中唯一的 `sum` 是四个 `MemXfer*` 项。其他所有内容都是 `max`/blend。任何把计算 lane 串行化的重新实现都会立即偏离二进制。
### 示例
一个 bundle 包含以下项（所有周期整数都取 [Per-Opcode Cycle Constants](per-opcode-cycle-constants.md) 的 `6acc60406` 列）：一个 bf16 matmul/latch issue（class `0x05`，向 `R[1]` 写入 212），一个 primary matrix-result read（class `0x1b`，向 `R[2]` 写入 127），以及一个 input DMA（向 `R[9]` 写入 30 startup + 向 `R[10]` 写入 64 bandwidth；DMA 项只是说明性示例，并非字节锚定）。
```text
R[1]=212 R[2]=127 R[9]=30 R[10]=64 (other slots 0)
mem = R[9]+R[10]+R[11]+R[12] = 30 + 64 + 0 + 0 = 94 (group B serial sum)
acc = max( vec_alu=0, mem=94 ) = 94
acc = max( 94, R[0]=0, R[1]=212, R[2]=127, ... ) = 212 (group C plain max)
bundle cost = 212 cyclesmatmul lane 是瓶颈；matrix-result read 和 DMA 免费重叠在它下面。这就是该模型的全部要点。
Bundle 组装 — CostModel::GetCycles
CostModel::GetCycles @ 0x130aade0 是单条 HLO 指令的 per-op resources 被收集并组合进调用者将要归约的累加器的位置。它是调度器按节点调用的 bundle 级入口。
// xla::jellyfish::CostModel::GetCycles(hlo, opts, *out_rv) @ 0x130aade0 (decompiled, paraphrased)
ResourceVector tmp = {}; // 184-byte zeroed RV (v31)
Status s = CostModel::GetHloResourcesImpl(&tmp, hlo, opts, ...); // fill slots via Acc/Record*
if (!s.ok()) return s; // AddSourceLocationImpl(cost_model.cc:2025)
double scalar = tmp.scalar_cycles; // var_220 → var_430, the cross-bundle scalar
AddOptions ao = ResourceVector::AddOptions::Defaults();
ResourceVector::Add(out_rv, &tmp, &ao); // combine tmp into the running accumulator
out_rv->scalar = scalar; // [rbx+8]
```text
`GetHloResourcesImpl` 会逐 slot 填充 `tmp`（上面的 `Acc`/`Record*` 写入家族）。随后 `GetCycles` 用默认 `AddOptions` 调用 `Add`，把 `tmp` 合入调用者的累加器。`[rbx+8]` 处的标量携带非向量化周期项（例如固定 transcendental 估计）并与向量并行存在。调用者之后对累加后的向量调用 `MaxResourceCycles` 并加上标量；显式的 `MaxResourceCycles + scalar` 序列见下面的 loop emitter。
---
## 跨向量组合 — `Add`
当两个区域被打包在一起（融合计算中的两个 op，两个 sub-bundle）时，它们的成本向量由 `Add` @ `0x1c89b820` 组合。规则是**几乎所有 slot 按 slot 求和，但两个 DMA-latency slot 取 `max`**：
```c
// xla::jellyfish::ResourceVector::Add(this, other, AddOptions) @ 0x1c89b820 (decompiled, paraphrased)
// most slots: this[i] += other[i]
// slots 9 and 11: if (AddOptions.flag == 1) this[i] = max(this[i], other[i]) else +=
for (int i = 0; i < 23; ++i) {
bool is_dma_latency = (i & 0x1D) == 9; // selects exactly i==9 and i==11
if (is_dma_latency && AddOptions.flag == 1)
this[i] = max(this[i], other[i]); // DMA startup paid once, not twice
else
this[i] += other[i]; // bandwidth + compute accumulate
}
AddToOperandBytesAccessed(this, other); // merge per-operand byte maps (+0xb8)
AddToOutputBytesAccessed(this, other); // merge per-output byte maps (+0xd8)掩码 (i & 0x1D) == 9 是二进制用于选择索引 9 和 11 的紧凑写法：9 & 0x1D == 9 且 11 & 0x1D == 9，而 [0,22] 中其他索引都会映射到别处。这两个索引是 MemXferInputLatency（R[9]）和 MemXferOutputLatency（R[11]），即 DMA startup 项。反编译中的前导 SIMD 块（在 [+0x40]/[+0x50] 上的 vcmpnltpd + vmaskmovpd）正是对 R[8..11] 窗口执行同一个“latency slot 取 max、bandwidth slot 相加”决策的向量化形式；标量尾循环处理其余部分。
NOTE — 摊销 DMA 模型。 组合两个 sub-vector 不应重复计算 DMA 的固定启动延迟：如果两个子区域从同一个输入流读取，引擎只支付一次 startup，但支付两次 bandwidth。
Add正是这样编码的：两个 latency 项取max，两个 bandwidth 项（R[10]、R[12]）以及所有计算项求sum。AddOptions.flag（其Defaults为AddOptions::Defaults@0x1c89ada0）切换这一行为；关闭该 flag 时，latency 项也会sum。
循环缩放 — ScaleResource 和 GetSubset
软件流水化循环体会定价一次并乘以 trip count，但 DMA startup 不应缩放。ScaleResource @ 0x1c89b6a0 用一个标量乘以一个 slot，跳过两个 DMA-latency slot：
// xla::jellyfish::ResourceVector::ScaleResource(Resource r, double factor, ScaleOptions) @ 0x1c89b6a0
if ((r & 0xFFFFFFFD) == 9 && !ScaleOptions.flag) // r==9 or r==11 (DMA latency)
return; // startup is paid once → do not scale
if (r >= 0x17) trap(); // bound: 23 slots
this[r] *= factor; // bandwidth/compute scale linearly
// r==10 / r==12: also scale the per-operand / per-output byte-map entries (flat_hash_map walk)
```text
`(r & 0xFFFFFFFD) == 9` 同样是“这是否为 DMA-latency slot”的测试（清除 bit 1，因此 `9` 和 `11` 都匹配 `9`）。当缩放 *bandwidth* slot（`r==10`、`r==12`）时，函数还会遍历 per-operand / per-output bytes-accessed 的 `flat_hash_map` 并缩放每个条目，因为字节量随 trip count 缩放。
`GetSubset` @ `0x1c89bb00` 把每迭代向量划分为 prologue / steady-state / tail 子集，由四个 boolean `SubsetOptions` flag（`a3[0..3]`）控制。它先把目标清零，然后按相同的 `(idx & 0x1D) != 9` 在某些 flag 组合中排除 DMA-latency 项，并在条件满足时对每个 source slot 做 `+=` 累加：
```c
// xla::jellyfish::ResourceVector::GetSubset(SubsetOptions) @ 0x1c89bb00 (decompiled, paraphrased)
// dest := 0
// for each slot i in 0..22, gated by the four flag bits:
// if slot selected: dest[i] += src[i]
// the per-flag predicates exclude the DMA-latency slots (i==9, i==11) via (i & 0x1D) != 9
// flag[1] also merges operand-bytes (+0xb8); flag[3] also merges output-bytes (+0xd8)结合 ScaleResource 后，净效果是循环总成本为 prologue (input-DMA startup, paid once) + trip_count × (compute + DMA bandwidth) + tail (output-DMA, paid once)。MaxResourceCycles 把每个组合后的向量归约为标量 bundle 成本。
上下文中的归约 — 一个真实调用者
PartialReduceEmitter::EstimateCycles @ 0x10eac440 展示了端到端的完整模式，是 MaxResourceCycles 作为终端 bundle-cost 步骤的最清晰字节锚定证明：
// xla::jellyfish::PartialReduceEmitter::EstimateCycles(op_windows) @ 0x10eac440 (decompiled, paraphrased)
ResourceVector rv = {}; // local 23-slot vector (v45..)
int64 compute = 0;
for (each input window) {
compute += per_window_cycles * Product(window_shape); // scalar compute term (v62)
}
CycleTable *ct = CycleTable::Create(target);
for (each operand) { // deposit memory cost into rv slots
RecordInputMemXferCycles(rv, param, ...); // → R[9], R[10]
RecordOutputMemXferCycles(rv, param, ...); // → R[11], R[12]
}
double bundle = ResourceVector::MaxResourceCycles(&rv); // reduce 23 slots → scalar
int64 cyc = (int64)bundle; // vcvttsd2si
result->cycles = cyc + compute; // bundle bottleneck + compute term
```text
这些写入填充 memory slot；`MaxResourceCycles` 将向量折叠为瓶颈 lane；结果是 `(int)max-lane + compute`。整数转换是 `MaxResourceCycles` 调用后紧跟的字面 `vcvttsd2si rax, xmm0`，加法（`rax + v62`）写入 emitter 的 cycle 字段。同样的 `MaxResourceCycles → (int) → + scalar` 形态也出现在 `BaseCostModelMetricCalculator::Calculate`（`0x1304ee00`）、`CostModel::GetLoopFusionOrUnfusedHloCycles`（`0x130b2bc0`）、`CostModel::GetCyclesIfFused`（`0x130aba40`）、`CostModel::GetOutputFusionOrConvolutionCycles`（`0x130aede0`）和 `ParamInput::EstimateComputeCycles`（`0x1126ee20`）中。唯一不是成本 emitter 的调用者是 SDC sequence-checker 的 `MaybeInjectMxuSequences::$_3`（`0x144fc5c0`），它调用 `MaxResourceCycles` 在验证注入的 MXU sequences 时重新推导候选 bundle 的 issue cost；它与上面的 `AccumulateInstructionUsage` 写入 lambda 位于同一个 translation unit。这七个函数是二进制中 `MaxResourceCycles` 的完整调用者集合。
---
## 依赖延迟轴 — `LatencyBetween`
Bundle 吞吐成本（上文）是调度器的两个输入之一。另一个是两个 op 之间的读后写依赖延迟，即调度器在 consumer 读取 producer 结果前必须插入的最小 cycle gap。这是一个*不同*的表家族（[CycleTable Family](cycletable-family.md) 记录了这个拆分），且按每个 op-*pair* 组合，而不是按每个 bundle。
`LatencyTable::LatencyBetween(LloValue, LloValue)` @ `0x1c89f820` 是公开 dispatcher。它调用每代的 `LatencyBetweenInternal` virtual，然后应用两个跨代一致的修正：
```c
// xla::jellyfish::LatencyTable::LatencyBetween(a, b) @ 0x1c89f820 (decompiled, paraphrased)
int lat = this->vtable[+0x18](this); // base: per-gen LatencyBetweenInternal
if (this->jitter_bitgen) // optional: +Uniform[0,101) jitter (+0x10)
lat += UniformInt(0, 101);
uint16 oa = a->opcode, ob = b->opcode;
if (oa == 132 && ob == 132) // matmul→matmul: floor from AutoOr<int>
lat = max(lat, autoOr.value_or(16));
else if (oa == 130 && (ob - 130) <= 2 && lat < 3) // matprep→matres family: floor 2
lat = 2;
return lat; // VLOG(2): "resolved latency from <a> to <b>: N"每代的 LatencyBetweenInternal 本身是在每代 latency array 的相关字段上，按每个 op-pair 取 max。Jellyfish 形式 LatencyTableJellyfish::LatencyBetweenInternal @ 0x1c8a0d60 按字节精确读取为：从 v15 = 1 开始，然后对每个匹配的 producer/consumer opcode 关系，把 v15 提升到对应 latency-array 字段（store-then-load 链用 v5[6]，pseudo-EUP 用 v5[7]，RPU/result-FIFO 类用 v5[8]/v5[11]/v5[19]/v5[20]，按 latch-mode 高位键控的 matprep/matmul/transcendental 类用 v5[9..18]），并有硬性下限 4（indexed-store-followed-by-load 和 set-IAR-followed-by-indexed-load 会提升到 5）。整个组合运算都是 if (v15 <= field) v15 = field;，也就是 max。因此两个 op 之间的依赖延迟是所有适用 RAW-hazard 字段的最大值，精确镜像 bundle-cost 的 max。
NOTE — 两个
max，两个轴。MaxResourceCycles在一个 bundle 内的功能单元 lane 上取 max（吞吐/占用）。LatencyBetween在一个 op pair 的 hazard 字段上取 max（依赖/停顿）。两者都是max，两者都来自同一个每代Performance对象（Performance Family Overview），但它们回答的是不同问题，调度器会同时消费两者。
送入调度器
CostModelLatencyEstimator::GetLatencyBetween(HloGraphNode, HloGraphNode) @ 0x10ff8f00 是 LatencyHidingScheduler 查询的 HLO 图级 estimator。它在 flat_hash_map<HloInstruction*, double>（estimator +0x18/+248）中 memoize 结果；对带计算的 opcode pair，它通过 CostModel::GetCycles（上面的 bundle 组装路径）再经过 per-bundle MaxResourceCycles 归约。对 async / collective 边，它改用 collective overlap 模型，包括 PCIe 和 DCN 带宽项（available_pcie_bandwidth、kDcnEstimatedBandwidth、kPcieMinimumTime、kDcnMinimumTime，都在首次调用时通过 __cxa_guard 懒计算一次）以及 SparseCore custom-call 估计（GetSparseCoreCustomCallCostEstimateLatency、GetRadixSortEstimatedExecutionCycles、gather/scatter offload estimators）。它返回的 cycle count 会通过 Target::TensorCoreFrequencyInMegaHertz 转为 wall-clock seconds（kCyclesPerMicrosecond guard variable，等于 TC clock in MHz；见 Cost Model Overview）。
随后调度器使用这两个量：每节点 bundle 成本（本页的 MaxResourceCycles）和每条边依赖延迟（LatencyBetween），来计算 list-scheduling priority。当依赖延迟允许时，它会把 async work 重叠在 compute 下面，并把独立 op 打包到 bundle 中；bundle 成本是 max-lane occupancy，而不是 sum。调度器内部（list priority、async tracking、独立的 47-ID ResourceType concurrency model）见 Scheduler Overview 和 ResourceType Taxonomy 页面。
GOTCHA — Bundle 成本是 per-lane
max，绝不是全局 sum。整个流水线中唯一的加法是归约内部四项 memory group 以及 packed vectors 之间的 per-slotAdd（它本身又对两个 DMA-startup 项做max组合）。把 bundle 定价为 slot 成本之和会给每个 multi-lane bundle 定错价。
交叉引用
- Resource Enum — 23-slot
ResourceVector、Acc写入，以及本页消费的MaxResourceCycles三组归约介绍。 - CycleTable Family — 向每 op 写入提供输入的
GetResource（op→slot）和GetCyclesForThroughput（per-class throughput）；吞吐与延迟的拆分。 - Per-Opcode Cycle Constants — 归约前写入 slot 的每代 cycle 整数。
- Cost Model Overview — 三个类家族、
Targetclock-frequency wiring，以及每个 cycle number 背后的 factory dispatch。 - Performance Family Overview — 吞吐轴和延迟轴共同读取的每代
Performancegrid。 - MXU Latency Overview — 为 matmul/matprep 写入提供的 per-
(MatmulModifier × Resource)reservation matrices。 - ResourceType Taxonomy — 调度器独立的 47-ID concurrency-cap enum（不是成本模型的
Resource）。 - Scheduler Overview — 消费 bundle cost 和 dependency latency 的
LatencyHidingScheduler。 - Bundle Model Overview — 本页定价其 issue cost 的 VLIW bundle 字。
- Learned Cost-Model Client — 为什么这里的每个数字都是静态
.rodata常量，而不是 ML prediction。 - Binary:
extracted/libtpu-0.0.40-cp314-cp314-manylinux_2_31_x86_64/libtpu/libtpu.so（build-id89edbbe81c5b328a958fe628a9f2207d） - Index entry: Part VII — Cost & Latency Model / Core model — 返回索引