NormalizedComputationCost

本页中的所有地址都适用于 libtpu-0.0.40-cp314 wheel 中的 libtpu.so（BuildID md5 89edbbe81c5b328a958fe628a9f2207d，781,691,048 字节，未剥离符号 — 每个符号都是反混淆后的 C++ 名称）。段映射：.text/.rodata VMA == 文件偏移；.data.rel.ro VMA − 0x200000 == 文件偏移。所有地址均为 VMA。其他 libtpu 构建会有所不同。

摘要

TpuPriorityFusionQueue::NormalizedComputationCost（@ 0x130989a0）是 TPU priority-fusion 队列在为 producer→consumer 融合排序时扣除的计算项。它针对每个 HLO op 回答一个问题 — 这个 op 会给融合区域增加多少向量/矩阵工作量 — 并返回一个 double。不同于 bundle-occupancy 模型（TpuHloCostAnalysis）会把每个 op 的周期存入 23 槽 ResourceVector，此函数会把每个 op 折叠成一个标量权重，并乘以 Target::ChunksIn(shape)（该 op 的 chunk 粒度元素数）。LLVM 中的类比是 TargetTransformInfo::getInstructionCost 查询，它返回粗粒度的“开销类别”，而不是周期精确的延迟 — 并且像 TTI 一样，它是少数离散权重构成的阶梯，而不是连续模型。

权重阶梯是一个针对 HloOpcode 的 switch，编译为 .rodata 0xae0dcdc 处包含 106 项的自相对跳转表（index = opcode − 0x18，opcode > 0x81 → default）。它正好解析为六个标量层级 — 0.0（免费的数据布局 op）、1.0（默认的廉价 elementwise）、2.0（融合 parameter 读取）、4.0（logistic / reduce / 跨 lane broadcast）、10.0（divide）、42.0（erf）— 外加三个结构性逃逸：convolution 不走阶梯，而是按 flop_count / peak 计价；fusion 通过递归求和其 body 的逐 op 成本计价；dot 是 CHECK 致命错误，因为 dot 在到达这里之前必须已被 lowered 为 convolution。

GetCyclesIfFused（@ 0x130aba40）是优先级公式的另一半 — 融合 (producer, consumer) 对的 bundle 成本。它不是手写的 producer+consumer 合并；它构建一个 FusionState，用 producer 扩展 consumer 的 operand 集合，然后通过用于任意 fusion 的同一套 GetHloResourcesImpl 机制为合并后的 op 重新计价。跨 functional-unit 的“打包” — 连续 MXU op 重叠、producer→consumer HBM 往返被丢弃、input-DMA 启动延迟只付一次 — 发生在该共享机制内部（ScaleAndSumOutputFusionResourceVectors @ 0x130b8320），这就是 total_fused < total_unfused 且融合优先级为正的原因。本页记录两者：opcode→weight 表及其字节精确常量，以及 GetCyclesIfFused 的融合合并驱动。

对于重新实现，契约是：

• 六层 opcode→weight 表、.rodata 权重常量，以及每一类的理由（为什么 divide=10、erf=42，为什么 reduce 使用 operand 的 ChunksIn，以及 parameter 只对前两个槽位为 ×2）。
• LoopFusion operand 累加前置路径及其元素类型门控掩码。
• broadcast 跨 lane 路径和 convolution flop→cycle 公式。
• fusion 递归求和并缓存路径（它不是纯查表）。
• GetCyclesIfFused 的资格门控、conv-like 主 op 选择、max-pool FLT_MAX 哨兵，以及 FusionState → GetHloResourcesImpl → ResourceVector::Add 合并驱动。
• 四个 sub-emitter 的 ScaleAndSumOutputFusionResourceVectors 组合，以及显式的 slot-9/11 MAX 组合。

计算权重	TpuPriorityFusionQueue::NormalizedComputationCost(HloInstruction*, long) @ 0x130989a0
跳转表	.rodata 0xae0dcdc — 106 × i32 自相对（index = opcode − 0x18）
权重层级	0.0, 1.0, 2.0, 4.0, 10.0, 42.0 + conv-flop / fusion-recurse / dot-fatal
Chunk 计数	Target::ChunksIn(Shape&) @ 0x1d619900
融合合并	CostModel::GetCyclesIfFused(producer, consumer, opts, ResourceVector*) @ 0x130aba40
合并核心	ScaleAndSumOutputFusionResourceVectors @ 0x130b8320（4 个 sub-emitter；slots 9/11 MAX）
资格门控	IsFusionSupportedHlo @ 0x130abee0
源文件	tpu_instruction_fusion.cc（weight） · cost_model/cost_model.cc（fused merge）

---

NormalizedComputationCost — Opcode→Weight 阶梯

目的

这是 bundle 感知融合优先级（priority = total_unfused − total_fused）中的计算惩罚部分。对于一个候选融合边，它估计 consumer（或递归地，一个嵌套 fusion）增加了多少向量/矩阵工作量。结果是一个 double；优先级队列会按用户数量缩放后扣除它。函数接受 (HloInstruction* inst, long operand_index)：rbx 是所属的 TpuPriorityFusionQueue this（持有 conv flop 缓存，以及位于 this+0x170 / *((q*)this+46) 的 Target*）；inst 是要计价的 op；operand_index 同时驱动递归保护和 parameter 槽位层级。

入口点

NormalizedComputationCost (0x130989a0)                 ── double, computes a scalar weight
  ├─ (operand_index>0 + binary + iota/broadcast op0)   ── recursion guard → flop/return tail
  ├─ LoopFusion pre-path (0x13098a13)                  ── multi-operand elementwise estimate
  │     └─ Target::ChunksIn (0x1d619900) × (1 + #non-zero-minor operands)
  └─ per-opcode switch (0x13098b62)                    ── jump table @ .rodata 0xae0dcdc
        ├─ 0.0  block  0x130995ee   (layout/metadata ops)
        ├─ 1.0  block  0x13098e2f   (DEFAULT — cheap elementwise)
        ├─ 2.0  block  0x13098de1   (parameter, operand_index ≤ 1)
        ├─ 4.0  blocks 0x13098d8e / 0x13098d74 / 0x13098b87  (logistic / reduce / broadcast)
        ├─ 10.0 block  0x13098d52   (divide)
        ├─ 42.0 block  0x13098e49   (erf)
        ├─ conv block  0x13098db0   (flop_count / peak — see CONV FORMULA)
        ├─ fusion block 0x13098e09  (cached, else recursive sum of body)
        └─ dot block   0x130996c5   (CHECK-fatal "Dots should have been replaced…")
```text

### 算法

```c
double NormalizedComputationCost(this /*rbx*/, inst /*r14*/, operand_index /*r15*/):  // 0x130989a0
    // (0) RECURSION GUARD — when called per fused-edge with operand_index>0,
    //     a binary op whose operand(0) is iota(0x43) or broadcast(0x1a) returns 0.0.
    if operand_index > 0:
        if (inst.operand_count_field & ~1) == 2:          // [inst+0x10] low form == 2
            op0 = inst.operand(0).opcode;                 // [operand(0)+12]
            if op0 == 0x43 /*iota*/ or op0 == 0x1a /*broadcast*/:
                return 0.0;
        goto per_opcode_switch;                           // LABEL_29

    // (1) LOOP-FUSION PRE-PATH — multi-operand elementwise estimate.
    if operand_index == 0
       and inst.IsLoopFusion()
       and root_element_type in NUMERIC_MASK            // 0x2FFF91FFE | {0x20,0x21} | 0x400048000
       and inst.fused_instructions_computation().instruction_count() <= 254:  // [comp+88]
        multiplier = 1.0;                                 // qword_A2DF230
        if inst.operand_count >= 2:                       // [inst+0x10] >= 2
            for each operand i:
                op_minor   = operand(i).shape.dimensions()[0] >> 1;   // minor dim, halved
                if op_minor != 0: multiplier += 1.0;      // +1 per NON-zero-minor operand (gated)
                root_minor = root.shape.dimensions()[0] >> 1;
                if op_minor >= root_minor: break;         // operand as-wide-or-wider → abandon
            else:
                return Target::ChunksIn(root.shape) * multiplier;     // estimate kept
        // any operand as-wide-or-wider, or 0/1 operands → fall through to the switch
        goto per_opcode_switch;

per_opcode_switch:                                        // 0x13098b62
    xmm0 = 0.0;                                           // pre-zeroed (the 0.0 tier needs no store)
    op = inst.opcode;                                     // byte [inst+0xc]
    if (op - 0x18) > 0x69: goto default_1_0;
    switch (op):                                          // jump table @ .rodata 0xae0dcdc

      case 0x18,0x27,0x29,0x2A,0x43,0x61,0x81:            // bitcast/concat/constant/convert/
        return 0.0;                                       //   iota/reshape/tuple — no compute

      case 0x1A /*broadcast*/:  return BroadcastWeight(inst);          // see BROADCAST PATH
      case 0x2B /*convolution*/: return ConvolutionWeight(this, inst); // see CONV FORMULA
      case 0x32 /*divide*/:     return Target::ChunksIn(root) * 10.0;  // qword_A2DF498
      case 0x38 /*erf*/:        return Target::ChunksIn(root) * 42.0;  // qword_A2DF1A0
      case 0x47 /*logistic*/:   return Target::ChunksIn(root) *  4.0;  // qword_A2DE830
      case 0x5B /*reduce*/:     return Target::ChunksIn(operand(0).shape) * 4.0;  // OPERAND, ×4
      case 0x52 /*parameter*/:
        if operand_index <= 1:  return Target::ChunksIn(root) * 2.0;   // vaddsd self
        else:                   return 0.0;
      case 0x3D /*fusion*/:     return FusionWeight(this, inst);       // see FUSION PATH
      case 0x34 /*dot*/:        LOG(FATAL) "Dots should have been replaced by convolutions.";
                                                          // tpu_instruction_fusion.cc:863

      default:                  return Target::ChunksIn(root) * 1.0;   // vcvtsi2sd, no vmulsd

注意 — 0.0 层级不需要额外工作，因为分发前 xmm0 已被清零（0x13098b5e 处的 vxorpd xmm0,xmm0）；1.0 默认情况只发出对 ChunksIn 的 vcvtsi2sd，没有乘法。其他所有层级都是用一个 .rodata double 执行 vmulsd。重新实现可以把 0.0 和 1.0 情况折叠进 chunk-count 计算中，但 dot 情况是硬 CHECK — 它必须中止，而不是返回一个数。

权重表

root 是 inst.shape；ChunksIn(s) = Target::ChunksIn(s)（@ 0x1d619900）是该 op 的 chunk 粒度元素数。跳转表目标已从 .rodata 0xae0dcdc 字节验证；权重常量已从 .rodata 字节验证。

权重 / 路径	常量	块 @	Opcodes（hex = name）
0.0 — 无成本	（xmm0 预清零，ret）	0x130995ee	0x18 bitcast, 0x27 concatenate, 0x29 constant, 0x2A convert, 0x43 iota, 0x61 reshape, 0x81 tuple
1.0 — ChunksIn(root)（DEFAULT）	（无 vmulsd）	0x13098e2f	下方未列出的每个 opcode 0x19..0x80（廉价 unary/binary elementwise、structural、collective、control-flow、I/O）
2.0 — ChunksIn(root) × 2	vaddsd self	0x13098de1	0x52 parameter — 仅当 operand_index ≤ 1；否则 0.0
4.0 — ChunksIn(root) × 4	0xa2de830 = 4.0	0x13098d8e	0x47 logistic
4.0 — ChunksIn(operand(0)) × 4	0xa2de830 = 4.0	0x13098d74	0x5B reduce（按被 reduce 的 operand 计价，而不是 root）
4.0 — broadcast cross-lane	0xa2de830 = 4.0	0x13098b87	0x1A broadcast（条件性 — 见 BROADCAST PATH）
10.0 — ChunksIn(root) × 10	0xa2df498 = 10.0	0x13098d52	0x32 divide
42.0 — ChunksIn(root) × 42	0xa2df1a0 = 42.0	0x13098e49	0x38 erf
conv flop 路径	（见 CONV FORMULA）	0x13098db0	0x2B convolution
fusion 递归 / 缓存	flat_hash_map @ this+0x80	0x13098e09	0x3D fusion
dot CHECK-fatal	LogMessageFatal	0x130996c5	0x34 dot

跳转表正好有 11 个不同目标：上面列出的十个加上共享的 0.0 块（0x130995ee）。Opcode→name 解析遵循按字母顺序排列的 XLA HloOpcode 枚举（0x18 kBitcast … 0x81 kTuple）；函数使用的加载期值 — 0x34 dot（fatal）、0x38 erf、0x47 logistic、0x52 parameter、0x5B reduce — 由函数自身的 opcode 比较相互印证（例如 GetCyclesIfFused 针对 reduce 特例测试 opcode != 91，91 == 0x5B）。

考量 — 为什么阶梯有这些层级

0.0/1.0 之上的四个标量层级排列的是 elementwise 开销，而不是 op 类别：

• 0.0 op 是纯数据布局 / 元数据 — bitcast、reshape、convert、concatenate、constant、iota、tuple。融合它们不会增加 functional-unit 工作量，因此它们从不惩罚优先级。这与 bundle 模型的 ZERO 分支一致（TpuHloCostAnalysis）。
• divide 是 10.0 — 它是最昂贵的廉价 elementwise op：一个 reciprocal-plus-multiply 微序列。bundle 模型也一致，会把 divide 展开为四次存入的 reciprocal 序列。
• erf 是 42.0 — 模型中单个最昂贵的 elementwise op：一段长 polynomial。这个 42.0 是 elementwise-erf 权重，不是 matmul/conv 类别。Matmul 和 convolution 由下方独立的 flop 路径计价；不存在 ×42 matmul 分支。
• reduce 独特地乘以 OPERAND 的 ChunksIn（operand(0).shape），而不是 root 的。reduce 的计算随其 reduce 的大输入缩放，而不是随较小的 reduced 输出缩放 — 这是 HandleReduce 的 ExtentProduct(operand) flop 公式在优先级中的镜像。
• parameter 只有在 operand_index ≤ 1 时为 ×2 — 前两个 fused-parameter 槽位承担 param tensor 的 read-plus-forward；更高槽位免费。这是唯一由 operand_index 而非 opcode 门控的层级。

易错点 — 42.0 是 erf 的 elementwise 权重（块 0x13098e49，常量 0xa2df1a0），不是 matmul/conv 类别。Convolution 走独立的 flop_count / peak 路径（块 0x13098db0）；没有 opcode 映射到 ×42 matmul 层级。把 matmul 按 ×42·ChunksIn 计价，在大型 conv 上会错几个数量级。

---

LoopFusion Operand 累加前置路径

目的

在逐 opcode switch 之前，loop（elementwise）fusion 会在 0x13098a13 获得一个特殊的 multi-operand 估计。直觉是：loop fusion 的计算量大约等于其输出 chunk 数，乘以每个必须 broadcast/expand 到输出宽度的 operand 的一个单位。一个已经和输出一样宽的 operand 不增加额外成本；较窄的 operand（需要 lane/sublane 扩展）增加一个单位。

算法

// reached only when operand_index == 0
if inst.IsLoopFusion()
   and element_type_in_numeric_mask(root.shape.element_type)
   and fused_instructions_computation().instruction_count() <= 254:   // 0xFE cap, [comp+88]
    multiplier = 1.0;
    if inst.operand_count >= 2:
        for each operand i in [0 .. operand_count):
            op_minor   = operand(i).shape.dimensions()[0] >> 1;     // [operand.shape+8] >> 1
            if op_minor != 0: multiplier += 1.0;                    // gated: skip zero-minor operands
            root_minor = root.shape.dimensions()[0] >> 1;           // re-read [inst.shape+8] each iter
            if op_minor >= root_minor:
                goto per_opcode_switch;          // operand as-wide-or-wider → abandon estimate
        return Target::ChunksIn(root.shape) * multiplier;
    goto per_opcode_switch;                       // 0 or 1 operand → switch on the root opcode
```text

数值元素类型门控与成本模型中处处使用的掩码相同：对 `et ≤ 0x21` 使用 `_bittest64(0x2FFF91FFE, et)`，或 `(et & ~1) == 0x20`，或对 `et ≤ 0x22` 使用 `_bittest64(0x400048000, et)` — 覆盖 bundle 模型可以计价的数值与 packed 类型（同一掩码也出现在 `GetHloResourcesImpl` 和 `GetCyclesIfFused` 中）。`≤ 254` instruction-count 上限把估计限制在小 fusion 上。

> **易错点 —** 该循环对 operand 和 root 的 minor 维度都读取 `dimensions()[0] >> 1`，比较为 `op_minor >= root_minor`。`>> 1`（减半）对两侧对称应用，因此在比较中抵消 — 它是以 chunk 为单位的 *minor* 维度。一旦任意 operand 与输出一样宽，估计就会被放弃（落到 fusion-root opcode 上的 switch，即 case `0x3D`）；只有“所有 operand 都严格窄于输出”才保留估计。另有两个字节级细节：(1) `multiplier += 1.0` 增量**由 `op_minor != 0` 门控** — 一个减半后的 minor dim 为零的 operand 会被循环计入，但*不会*提高 multiplier（`0x13098…`/反编译第 732 行处 `v17 == 0` 时 `vaddsd` 结果被丢弃）；(2) 每次迭代都会从 `inst.shape` 重新读取 `root_minor`，而不是提升到循环外。因此保留的估计是 `ChunksIn(root) × (1 + #non-zero-minor operands)`，不是 `(1 + #operands)`。无条件返回估计，或为每个 operand 都提高 multiplier 的重新实现，会高估那些 operand 已经 full-width 或 zero-minor 的 fusion。

---

## Broadcast 路径（opcode `0x1A`，块 `0x13098b87`）

### 算法

```c
double BroadcastWeight(inst):                              // 0x13098b87
    target = this.Target;                                  // *((q*)this+46)
    if target[0x398] == 0: return 0.0;                     // broadcast-cost flag off (Target+920)
    operand = inst.operand(0);
    if operand.shape.rank > 3: return 0.0;                  // high-rank broadcasts not priced
    phys = LayoutUtil::MakeLogicalToPhysical(inst.shape.layout);
    sort(inst.dimensions());                                // ascending
    if ShapeUtil::IsEffectiveScalar(operand.shape):         return 0.0;   // splat of a scalar
    if phys.minor_most_dim is a broadcast dimension:        return 0.0;   // free minor-axis splat
    return Target::ChunksIn(inst.shape) * 4.0;             // cross-lane movement, qword_A2DE830

沿 minor-most（sublane/lane）轴的 broadcast 是免费的寄存器 splat；跨 lane 物化的 broadcast 成本为 4.0 中等权重。整条路径由 Target+0x398（= Target+920）处的每 Target broadcast-cost flag 门控 — 当该标志清零时，每个 broadcast 都是 0.0。

怪癖 — bundle-occupancy 模型无条件把 broadcast 计价为零（TpuHloCostAnalysis ZERO 分支），而此优先级模型可能对 cross-lane broadcast 收取 4.0·ChunksIn。这两个表面对布局 op 的处理有意不同 — bundle 路径把 broadcast 视为零 functional-unit occupancy；优先级路径则计入跨 lane 数据移动。不要把它们统一。

---

Convolution 公式（opcode 0x2B，块 0x13098db0）

算法

Convolution 完全逃离标量阶梯：它通过把 flop count 转换为 MXU cycles 来计价。flop count 按 unique_id 缓存，因此重复的优先级查询很便宜。

double ConvolutionWeight(this, inst):                      // 0x13098db0
    // (1) FLOP CACHE — flat_hash_map<int64,double> at this+0xd0 / this+208
    uid  = inst.unique_id();
    flop = this.flop_cache.find(uid);
    if MISS:
        TpuHloCostAnalysis ca(this.Target, /*model_tpu_specific=*/true);   // ctor 0x130a1620
        st = ca.HandleConvolution(inst);                   // 0x1e480be0
        CHECK(st.ok()) << "hlo_cost_analysis.HandleConvolution(instruction) is OK";
                                                            // tpu_instruction_fusion.cc:871
        flop = ca.flop_count(inst);                        // 0x1e4841e0
        this.flop_cache.emplace(uid, flop)                 // CHECK emplace.second @ :875
    // (2) FLOP → CYCLES
    if inst.batch_group_count() == 1 and inst.feature_group_count() == 1:   // dense conv
        fmt        = LhsFormatForConvInstruction(inst, this.Target);  // 0x1307bd40
        peak_flops = Target::FlopsPerSecond(fmt);          // vtable+0x718, 0x1d61f280
        freq_MHz   = Target::TensorCoreFrequencyInMegaHertz();        // 0x1d615b60
        peak_per_cycle = peak_flops / freq_MHz / 1e6;      // qword_A2E0208 = 1e6
        mxu_cycles = flop / peak_per_cycle;
        valu_slots = Target::VectorAluSlotsPerTensorCore();// vtable+0x500, 0x1d61e380 (int)
        derate     = Target[0x4ac] * (-0.03) + 1.0;        // A2E05A8 = -0.03; A2DF230 = 1.0
        return (valu_slots * mxu_cycles) / derate;
    else:                                                  // grouped / depthwise conv
        return flop * 0.00048828125;                       // A2E0118 = 1/2048
```text

dense conv 的成本是它的 flop count，按每 LHS-format 的峰值速率转换为 MXU cycles，再按 vector-ALU slot 数缩放，并除以 `(1 − 0.03·N)` headroom derate。Grouped convolution — MXU 对它们的加速不如 dense conv — 使用固定的 `flop/2048` 估计。`HandleConvolution` flop 与 [TpuHloCostAnalysis](tpu-hlo-cost-analysis.md#convolution--divides-by-both-group-counts) 中记录的是同一个（它会同时除以**两个** group count）；它遍历的 conv-shaped state 见 [ConvolutionCostState](convolution-cost-state.md)。

> **注意 —** 绝对的每代整数 `Target::FlopsPerSecond(format)`、`Target::VectorAluSlotsPerTensorCore()`，以及 `Target+0x4ac` derate operand 都来自每 codename / `chip_parts`，且**不**在此枚举（此构建只嵌入了 v7 芯片参数）。*公式*是字节精确的；常量由每代 MXU 页面拥有（[MXU Latency Overview](mxu-latency-overview.md)、[MatmulMode and Modifiers](matmul-mode-modifiers.md)）。`LhsFormatForConvInstruction` → `MatmulDataFormat` 选择也不在此解码（MEDIUM 说明每个 conv 选择哪个 peak）。

---

## Fusion 路径（opcode `0x3D`，块 `0x13098e09`）

### 算法

`fusion` op 按其 body 的逐 op 成本**求和**计价，递归执行，并按 `HloInstruction*` 缓存结果。这*不是*纯查表 — 缓存 miss 时，它遍历 fused computation 的 root chain，并对每条 fused instruction 递归调用 `NormalizedComputationCost`（推进 `operand_index`，这正是函数顶部递归保护用于防止 iota/broadcast operand 被重复计数的地方）。

```c
double FusionWeight(this, inst):                           // 0x13098e09
    // cache: flat_hash_map<HloInstruction*, double> at this+0x80 / this+128
    hit = this.fusion_cache.find(inst);                    // 0x13098e09
    if hit: return hit->second;

    comp = inst.fused_instructions_computation();          // 0x13098f1c
    cost = 0.0;
    // walk the fused instruction list (16-byte stride), skipping null slots,
    // tracking a (lo,hi) index pair into the computation's instruction array:
    for each fused_instruction f in comp:                  // LABEL_105..LABEL_112
        cost += NormalizedComputationCost(this, f, edge_index);   // recursion @0x13098f50
    this.fusion_cache.emplace(inst, cost);                 // SOO flat_hash_map insert
    return cost;

注意 — Case 0x3D 是一个缓存查找，miss（0x13098f18 之后）时会就地计算成本：它对 fused_instructions_computation() 的每条 instruction 递归调用 NormalizedComputationCost，累加进 xmm0（0x13098fd0 处的 vaddsd xmm0, var_30, xmm0），然后把和存入 this+0x80 map。缓存由 NormalizedComputationCost 自身在父 fusion 的第一次查询时填充，而不是由外部 pass 填充。

函数顶部的递归保护（operand_index > 0 + binary op + operand(0) 是 iota/broadcast → 0.0）防止 fused binary op 对一个已经由 loop pre-path 或 sibling edge 计入的 iota/broadcast operand 再次收费。

---

GetCyclesIfFused — 融合 Bundle 合并

目的

GetCyclesIfFused（@ 0x130aba40）返回融合 (producer, consumer) 对的 bundle 成本 — 优先级公式中的 total_fused 项。它的关键结构事实：它不会手工合并两个 ResourceVector。它把 consumer 当作一个 fusion，并用 producer 扩展其 operand 集合（一个 FusionState），然后用任意 fusion 都使用的同一个 GetHloResourcesImpl 为合并后的 op 计价。因此跨 FU packing 发生在共享机制内部，而不是此函数内部。签名是 StatusOr<double> GetCyclesIfFused(producer, consumer, const PerCalculationOptions&, ResourceVector* out)；rbx 是 StatusOr<double>* 返回槽，producer 通过 a4/v19 传入，consumer 通过 a3 传入，out 通过 a5/a6 传入。

入口点

GetCyclesIfFused (0x130aba40)                          ── StatusOr<double>
  ├─ IsFusionSupportedHlo      (0x130abee0)            ── eligibility gate → trivial 1.0 cy
  ├─ numeric element-type mask (0x2FFF91FFE | …)       ── non-numeric → 1.0 cy
  ├─ ShapeUtil::IsZeroElementArray                      ── (consumer ≠ reduce) zero-elem → 1.0 cy
  ├─ IsConvLowerable (0x14553620) / ExtractConvLikeHlo (0x1d6aa140)  ── conv-like main-op pick
  │     └─ GetReduceWindowType (0x1454d4a0) == -1/2     ── max-pool → FLT_MAX sentinel
  ├─ FusionState::Create(consumer, producer) (0x130ab320)           ── combined operand set
  ├─ GetHloResourcesImpl(consumer, opts, &fs, isFused=1) (0x130aa580) ── price merged op
  │     └─ ScaleAndSumOutputFusionResourceVectors (0x130b8320)       ── 4-emitter combine
  └─ ResourceVector::Add(out, merged, Defaults) (0x1c89b820)         ── fold into caller's out
```text

### 算法

```c
StatusOr<double> GetCyclesIfFused(producer /*a4*/, consumer /*a3*/, opts, out /*a5*/):  // 0x130aba40
    // (A) ELIGIBILITY GATES — return a trivial 1.0-cycle cost if the op is not modellable.
    if !IsFusionSupportedHlo(consumer, opts.target)        // 0x130abee0
       or consumer.shape.element_type not in NUMERIC_MASK  // 0x2FFF91FFE | {0x20,0x21} | 0x400048000
       or (consumer.opcode != 0x5B /*reduce*/              // reduce is exempt from the next test
           and ShapeUtil::IsZeroElementArray(producer.shape)):
        out.scalar = 0;                                    // *(this+1) = 0
        return StatusOr<double>(1.0);                      // *(this) = 1  (1 trivial cycle)

    // (B) CONV-LIKE MAIN-OP SELECTION — pick the conv/reduce-window "hero" between the two ops.
    main = consumer;                                       // default
    for cand in {producer, consumer}:                      // probe both
        if IsConvLowerable(cand):                          // 0x14553620
            conv = ExtractConvLikeHlo(cand);               // 0x1d6aa140
            if conv and conv.opcode == 0x5E /*reduce-window*/:
                t = GetReduceWindowType(conv);             // 0x1454d4a0
                // (C) MAX-POOL SENTINEL: a max-pool (t==2) or unknown (t==-1) window is not
                //     bundle-cost-modellable → emit FLT_MAX cycles so it never fuses.
                if t == 2 or t == -1:
                    out.resources[…] = FLT_MAX;            // 0x47EFFFFFE0000000 = qword_A2E0530
                    return StatusOr<double>(out.MaxResourceCycles());   // = FLT_MAX
            else:
                main = cand;                               // this cand is the conv-lowerable main op

    // (D) BUILD THE MERGED FUSION STATE — consumer augmented by producer's operands.
    FusionState fs;
    FusionState::Create(&fs, consumer, producer);          // 0x130ab320
    //   fs records consumer.operands() + producer.operands() as one combined set, and the
    //   per-producer operand-indices at which the consumer USES the producer (the internal edges).

    // (E) PRICE THE MERGED OP AS A FUSION.
    StatusOr<ResourceVector> merged =
        GetHloResourcesImpl(consumer, opts, &fs, /*isFused=*/1);   // 0x130aa580
    if !merged.ok(): return merged.status();               // propagate w/ source loc cost_model.cc:2154

    double fused_cycles = merged.value().scalar_cycles;    // [merged + 0x240-ish]

    // (F) FOLD THE MERGED BYTE/SLOT MAP INTO THE CALLER'S out VECTOR.
    out.Add(merged.value(), AddOptions::Defaults());       // 0x1c89b820 — slots 9/11 MAX, rest ADD
    return StatusOr<double>(fused_cycles);

GetHloResourcesImpl 通过 fs 遍历 consumer 的 fused expression 加上 producer。对于每条 operand edge，它会查询 IsProducerUse（@ 0x130ab0c0）：记录在 FusionState 中的 edge 是内部的，因此它们的 input-DMA 字节不会存入 MemXfer 槽 R[9..12] — 这就是 fused cost 中移除 producer→consumer HBM 往返的机制。producer 的 compute 会累加进与 consumer 相同的 ResourceVector，而 sub-emitter 成本由 ScaleAndSumOutputFusionResourceVectors 组合。

Max-Pool FLT_MAX 哨兵

一个值得指出的 early-exit：解析为 reduce-window（opcode 0x5E）且 GetReduceWindowType 为 2（max-pool）或 -1（unknown）的 conv-like op 不能用 bundle-cost 建模。函数不会错误计价，而是把 FLT_MAX 位模式 0x47EFFFFFE0000000（= .rodata 0xa2e0530，字节验证）写入结果 ResourceVector 的槽位，并返回其上的 MaxResourceCycles() — 一个实际无限的成本，保证该融合永远不会被选中。这与 GetHloResourcesImpl 路由中的同一个 t ∈ {−1, 2} reduce-window 哨兵相呼应（TpuHloCostAnalysis）。

---

ScaleAndSumOutputFusionResourceVectors — 每 FU 组合

目的

此例程（@ 0x130b8320）由 GetHloResourcesImpl 调用，用于把一个合并 op 的 sub-emitter 成本融合成一个 ResourceVector。它组合最多四个 sub-emitter 向量 — activations、kernel、output 和 conv-compute — 每个按自己的 iteration count 缩放，其中 memory-transfer latency 槽位（9 和 11）采用 MAX 组合，而不是求和。

算法

ResourceVector ScaleAndSumOutputFusionResourceVectors(            // 0x130b8320
        out, rv_act,  act_count,
             rv_kern, kern_count,
             rv_out,  out_count,
             rv_conv, conv_count):
    CHECK(conv_count >= act_count)  << "conv_compute_iteration_count >= activations_iteration_count";  // :1567
    CHECK(conv_count >= kern_count) << "… >= kernel_iteration_count";                                   // :1568
    CHECK(conv_count >= out_count)  << "… >= output_iteration_count";                                   // :1569
    out = 0;                                              // zero all 23 slots + byte maps

    // (1) ADD each scaled sub-emitter (default Add: every slot accumulates).
    for (rv, count, subset) in {(rv_act,  act_count,  SubsetOpts@0xae0f0c8),
                                (rv_kern, kern_count, SubsetOpts@0xae0f0cc),
                                (rv_out,  out_count,  SubsetOpts@0xae0f0d0),
                                (rv_conv, conv_count, SubsetOpts@0xae0f0d4)}:
        sub    = rv.GetSubset(subset);                    // 0x1c89bb00
        scaled = sub.GetScaled((double)count, ScaleOptions::Defaults());   // 0x1c89b3c0
        out.Add(scaled, AddOptions::Defaults());          // 0x1c89b820

    // (2) OVERRIDE slots 9 and 11 (MemXfer Input/Output LATENCY) with MAX-across-emitters × count.
    out.Acc(9,  max(rv_act[9],  rv_kern[9],  rv_out[9],  rv_conv[9])  * conv_count);   // +0x48
    out.Acc(11, max(rv_act[11], rv_kern[11], rv_out[11], rv_conv[11]) * conv_count);   // +0x58
    return out;
```text

四次 `GetSubset` 调用使用 `.rodata 0xae0f0c8 / 0xcc / 0xd0 / 0xd4` 处不同的 `SubsetOptions` 字面量（每个都是 4 字节 tuple `{0x00,0x01,0x01,0x01}`，字节验证）— 每个 emitter 一个 subset 形状。最后两次 `Acc` 调用是**显式 slot-9/11 MAX 组合**：input-DMA 和 output-DMA *启动延迟*项只付一次（取四个 sub-emitter 的最大值），而不是按 emitter 求和 — 传输的启动延迟会跨融合 sub-region 重叠，而其带宽（slots 10/12）和计算槽会累加。

> **注意 —** 此例程会*构建*组合后的 `ResourceVector`（其中 slots 9/11 采用 MAX 组合），并在 `Acc(this, 11, …)` 之后返回它；它自身**不会**约简为 cycles。最终标量约简发生在**调用方**，通过单独的 `MaxResourceCycles`（`@0x1c89b9e0`）完成。这里 slots 9/11 的 MAX 组合是显式的逐槽覆盖，不同于 `MaxResourceCycles` 约简对完整 `{9,10,11,12}` memory group 的串行求和（[Resource Enum](resource-enum.md#the-maxresourcecycles-reduction)）。

### 为什么 `total_fused < total_unfused`

```text
total_unfused = GetHloCycles(producer)·user_count + Σ_users GetHloCycles(user)
   each op is its OWN bundle: the producer's output is WRITTEN to HBM (output-DMA, R[11]/R[12])
   and each user READS it back (input-DMA, R[9]/R[10]). Those DMA cycles are counted (user_count+1)×.

total_fused = Σ_users GetCyclesIfFused(producer, user)
   the producer→user edge is INTERNAL (IsProducerUse drops its input-DMA), the producer's HBM
   output write is eliminated (it stays in VMEM), and the producer's compute slots OVERLAP the
   user's in the MaxResourceCycles plain-MAX group (e.g. producer Matmul R[1] overlaps user
   VectorAlu R[3..5]) instead of being two serial bundles.

priority = total_unfused − total_fused > 0  whenever the saved DMA + the bundle-packing overlap
           exceed any duplicated compute the fusion introduces.

对于单用户 producer，优先级是 (C_p + C_u) − C_f；对于 n 用户 producer，它是 (n·C_p + Σ_i C_u_i) − Σ_i GetCyclesIfFused(producer, user_i)，其中 C_p = MaxResourceCycles(RV_producer)，C_u 对 consumer 同理，C_f = GetCyclesIfFused(producer, consumer)。

---

IsFusionSupportedHlo — 资格门控

IsFusionSupportedHlo（@ 0x130abee0）是 GetCyclesIfFused 的第一个门控。当 consumer 不能被 bundle-cost 建模时，它返回 false（→ 平凡的 1-cycle 成本）：

• 元素类型具有 64 位 bit-width（ShapeUtil::ElementHasBitWidth(shape, 0x40)），除非该 op 是 custom fusion 或 collective-compute fusion；
• 该 op 是 zero-element array；
• opcode 位于 REJECT 集 _bittest(0x2000100000400001, opcode−5) 中，适用于 [5..0x42] 范围内的 opcode → {0x05 all-gather-done, 0x1B call, 0x31 custom-call, 0x42 infeed} — 即没有 functional-unit occupancy 的 control-flow / host-I/O op。

---

Worked Example — 一个小型 elementwise fusion

一个以 multiply（0x4B，默认层级 1.0）为 root 的 loop fusion，三个 operand，输出 [256,128]：

operand0 = parameter [256,128]   (root minor = 128)
operand1 = broadcast scalar      (minor << 128)
operand2 = parameter [256,128]   (minor = 128)
```text

LoopFusion pre-path：

```text
multiplier = 1.0
operand0: op_minor 128 >= root_minor 128  → break immediately (operand as-wide as root)
→ estimate abandoned → fall to per-opcode switch on the root (kFusion 0x3D) → recurse the body

对于 body 的 multiply leaf（默认层级），逐 op 权重为 ChunksIn([256,128]) × 1.0。把 root 替换为 divide，leaf 就是 ChunksIn × 10.0；替换为 erf，就是 ChunksIn × 42.0；替换为 dense convolution，就是 flop_count / peak（较重）。这是 total_unfused 要扣除的计算项；bundle 项 GetCyclesIfFused 通过上面的合并计算。

---

函数映射

函数	地址	作用
TpuPriorityFusionQueue::NormalizedComputationCost	0x130989a0	opcode→weight 标量 + conv/fusion 逃逸
Target::ChunksIn(Shape&)	0x1d619900	chunk 粒度元素数（×multiplier 基数）
TpuHloCostAnalysis ctor	0x130a1620	conv flop 子分析
HloCostAnalysis::HandleConvolution	0x1e480be0	conv flop emitter
HloCostAnalysis::flop_count	0x1e4841e0	读取缓存的 flop 属性
LhsFormatForConvInstruction	0x1307bd40	conv LHS → MatmulDataFormat（peak select）
Target::FlopsPerSecond	0x1d61f280	每 format 峰值（vtable+0x718）
Target::VectorAluSlotsPerTensorCore	0x1d61e380	VALU slot 数（vtable+0x500）
Target::TensorCoreFrequencyInMegaHertz	0x1d615b60	TC clock（cycles ← seconds）
CostModel::GetCyclesIfFused	0x130aba40	fused-pair bundle cost 驱动
IsFusionSupportedHlo	0x130abee0	资格门控（→ 1-cycle 平凡成本）
IsConvLowerable	0x14553620	conv-lowerable 谓词
ExtractConvLikeHlo	0x1d6aa140	拉取 conv/reduce-window root
GetReduceWindowType	0x1454d4a0	−1/2 max-pool 哨兵
FusionState::Create	0x130ab320	组合 operand 集 + internal-edge map
CostModel::IsProducerUse	0x130ab0c0	丢弃 internal-edge input DMA
CostModel::GetHloResourcesImpl	0x130aa580	为合并后的 op 计价
ScaleAndSumOutputFusionResourceVectors	0x130b8320	4-emitter 组合；slots 9/11 MAX
ResourceVector::Add	0x1c89b820	逐槽累加（Defaults）
ResourceVector::MaxResourceCycles	0x1c89b9e0	标量 bundle-cycle 约简

权重 / 公式常量（.rodata，字节验证）

地址	值	使用者
0xa2df230	1.0	默认权重 / conv derate +1.0 / multiplier 基数
0xa2de830	4.0	logistic、reduce、cross-lane broadcast
0xa2df498	10.0	divide
0xa2df1a0	42.0	erf
0xa2e0530	3.4028e38 (FLT_MAX)	max-pool GetCyclesIfFused 哨兵
0xa2e0208	1.0e6	conv freq_MHz → Hz
0xa2e05a8	-0.03	conv derate 斜率 1 − 0.03·Target[+0x4ac]
0xa2e0118	0.00048828125 (1/2048)	grouped-conv flop→cost 因子

---

相关组件

组件	关系
TpuHloCostAnalysis	提供本页缓存的 conv flop_count；标量阶梯的 bundle-occupancy 对应项
Resource Enum (23-slot)	ResourceVector 槽位和 fused merge 输入的 MaxResourceCycles 约简
ConvolutionCostState	HandleConvolution 在 flop 缓存前遍历的 conv-shaped state
Reduce-Window / Pooling Cost	max-pool FLT_MAX 哨兵背后的 GetReduceWindowType 分类
Per-Opcode Cycle Constants	存入合并后 ResourceVector 的每代 cycles

交叉引用

• TpuHloCostAnalysis — 提供 HandleConvolution / flop_count 的 flop/byte 模型，以及 bundle-occupancy RecordHloCycles 对应表面
• Cost Model Overview — 三类每代 class family，以及 conv 公式读取的 Target clock wiring
• Resource Enum (23-slot) — fused merge 约简所用的槽位名称、Acc 和 MaxResourceCycles 重叠模型
• ConvolutionCostState — flop 计价前构建的每 conv state
• Reduce-Window / Pooling Cost — reduce-window 类型分类和 max-pool 路径
• Per-Opcode Cycle Constants — 输入合并向量的每代 GetCyclesForThroughput 整数
• CycleTable Family — GetResource（op→slot）和 Instruction bucket enum
• MXU Latency Overview — 参数化 conv 公式的每代 FlopsPerSecond / VALU-slot 整数
• MatmulMode and Modifiers — 选择 conv peak rate 的 MatmulDataFormat 选择
• dot/conv MXU Lowering — 为什么原始 dot 到达此函数会是 CHECK-fatal
• Fusion Cost Model — 消费这些 compute 和 bundle cost 的优先级公式