GetHloResources 路由

地址适用于 libtpu-0.0.40-cp314 wheel 中的 libtpu.so（BuildID md5 89edbbe81c5b328a958fe628a9f2207d，781,691,048 字节，未 strip）。下面每个符号都是 demangled C++ 名称。.text/.rodata VMA == 文件偏移；.data.rel.ro VMA − 0x200000 == 文件偏移。所有地址都是 VMA。其他版本会不同。

摘要

CostModel::GetHloResourcesImpl 是 TPU bundle-occupancy 成本模型的路由主干：给定一条 HLO 指令，它决定由哪个成本子发射器为该 op 定价，并由该子发射器把该 op 的吞吐周期存入 23-slot ResourceVector。它是 flop/byte 侧的资源路由对应物。flop 侧（TpuHloCostAnalysis）遍历同一个 HLO，并写入供 fusion-priority 模型消费的 float 属性（+0x50 flops / +0x54 transcendentals / +0x58 bytes）；本页是另一个表面，即按 op 将周期存入 VLIW-bundle 资源向量。二者共享 HLO，并对哪些 op 昂贵达成一致，但它们由不同函数计算，存入不同结构，且从不共享任何数值。

路由分三个嵌套阶段发生，本页也围绕它们组织。**阶段 1 — GetHloResourcesImpl（@0x130aa580）**是五路顶层分派：collective 进入 network model；conv-lowerable op（或非 max-pool reduce-window）进入 MXU output-fusion model；collective-compute fusion 进入自己的模型；其余进入 loop/elementwise 默认路径。**阶段 2 — RecordCyclesIfFused（@0x130cc720）**是只从默认分支到达的 fusion-peel 路由器：对 fused region，它按 convolution、reduce-window、loop-fusion 和 output-fusion root 剥离到专用发射器，过程中丢弃 producer→consumer input-DMA 边，并对每个剩余 op 落到 leaf 发射器。**阶段 3 — RecordHloCycles（@0x130bbfe0）**是 leaf 按 op 存入：对 opcode 做 switch，编译为 .rodata 0xae0ebbc 处 0x7f 项自相对 jump table，把每个 opcode 路由到它发出的 CycleTable::Instruction 以及落入的具名 Resource 槽。

熟悉的参照框架是 LLVM TargetTransformInfo 成本路径，每个 opcode 到达一个 getInstructionCost 分支；不同之处在于这里的“cost”是存入结构化 bundle 向量的资源占用，稍后的 MaxResourceCycles reduction 会用显式重叠规则折叠该向量。本页记录三层分派、门控每个分支的谓词和 mask、leaf opcode→slot 路由，以及每个被路由的子模型如何写入 ResourceVector::Acc。由 TpuHloCostAnalysis 页面拥有的 flop-override 表不会在这里重复。

对于重新实现，契约是：

• 五路 GetHloResourcesImpl 分派：collective 谓词（带 opcode-61 collective-in-fusion 剥离）、numeric element-type bitmask 门、conv/reduce-window 选择以及 max-pool sentinel、collective-compute-fusion 与 loop/elementwise 默认路径。
• RecordCyclesIfFused 剥离顺序：conv → reduce-window → loop-fusion → output-fusion → leaf，以及在按 leaf 存入前运行的 IsProducerUse 边内部化。
• RecordHloCycles opcode→slot jump table：idx = opcode − 3 索引、ja > 0x7e → default 边界，以及每个具名 block 发出的 CycleTable::Instruction 和通过 CycleTable::GetResource 落入的 Resource 槽。
• leaf 存入机制：slot = GetResource(CT)，thru = cycletable.vtable[+0x10](CT)，Acc(slot, element_count × thru × W)。

Class	xla::jellyfish::CostModel (cost_model.cc)
阶段 1 — 顶层分派	CostModel::GetHloResourcesImpl @0x130aa580（5 路）
阶段 2 — fusion 剥离	CostModel::RecordCyclesIfFused @0x130cc720
阶段 3 — leaf 存入	CostModel::RecordHloCycles @0x130bbfe0
Leaf jump table	.rodata 0xae0ebbc（0x7f × i32 自相对；target = 0xae0ebbc + off）
Op→slot map	CycleTable::GetResource(Instruction) @0x1c89ce20 → dword_B438AEC[CT]
存入原语	ResourceVector::Acc(Resource, double) @0x1c89adc0（边界 < 0x17）
Public entry	CostModel::GetHloResources @0x130aa560（薄包装）
Source file	platforms/xla/service/jellyfish/cost_model/cost_model.cc

---

阶段 1 — GetHloResourcesImpl，五路顶层分派

目的

GetHloResourcesImpl 是入口：它为一条 HLO op 返回 StatusOr<ResourceVector>，方式是在五条定价路径中选择一条。它自身不会存入周期；每个分支都委托给构建并填充向量的子发射器。按求值顺序，分支是：collective → GetCollectiveCycles；conv-lowerable / 非 max-pool reduce-window → GetOutputFusionOrConvolutionCycles；collective-compute fusion → GetCollectiveComputeFusionCycles；默认 → GetLoopFusionOrUnfusedHloCycles（唯一会到达 RecordCyclesIfFused / RecordHloCycles 的分支）。

入口点

GetHloResources              @0x130aa560  ── public wrapper
  └─ GetHloResourcesImpl     @0x130aa580  ── StatusOr<ResourceVector>, 5-way
      ├─ IsSupportedCollectiveHlo  @0x130aeda0 ── arm 1 (+ opcode-61 fusion peel)
      │     → GetCollectiveCycles  @0x130abfc0     rv[+8] = scalar cycles
      ├─ IsFusionSupportedHlo      @0x130abee0 ── numeric-type + fusion-support gate
      ├─ IsConvLowerable           @0x14553620 ┐
      ├─ ExtractConvLikeHlo        @0x1d6aa140 ├ conv / reduce-window selection
      ├─ GetReduceWindowType       @0x1454d4a0 ┘  (type −1/2 max-pool → NOT conv-priced)
      │     → GetOutputFusionOrConvolutionCycles @0x130aede0   ── arm 2 (MXU)
      ├─ IsCollectiveComputeFusion @0x13e028c0
      │     → GetCollectiveComputeFusionCycles   @0x130b13a0   ── arm 3
      └─ (default)
            → GetLoopFusionOrUnfusedHloCycles    @0x130b2bc0   ── arm 4 → Stage 2
```text

### 算法

```c
StatusOr<ResourceVector> GetHloResourcesImpl(inst, opts, fs, isFused):  // sub_130AA580
    op = *(byte*)(inst + 12);                       // HloInstruction::opcode

    // (1) COLLECTIVE — all-reduce / all-gather / … priced by the network path.
    if IsSupportedCollectiveHlo(inst):              // sub_130AEDA0
        return GetCollectiveCycles(inst, &rv);      // sub_130ABFC0 — rv[+8] = scalar cycles
    if op == 61 (fusion):                            // peel a collective wrapped in a fusion
        comp = inst.fused_instructions_computation();
        c    = first fused instr where IsSupportedCollectiveHlo(c) holds;
        if c found and !IsCollectiveComputeFusion(inst):
            return GetCollectiveCycles(c, &rv);     // price the inner collective

    // (2) NUMERIC ELEMENT-TYPE + FUSION-SUPPORT GATE
    et = inst.shape.element_type;                   // *(u32*)*(inst+88)
    if (et <= 0x21 and bit(0x2FFF91FFE, et))         // numeric/packed-type mask (bittest64)
       or (et & 0xFFFFFFFE) == 0x20                  // (the F8 pair 0x20/0x21)
       or (et <= 0x22 and bit(0x400048000, et)):     // the loop/output-fusion peel mask
        if IsFusionSupportedHlo(inst, opts.target):  // sub_130ABEE0
            // (3) CONV / REDUCE-WINDOW MAIN-OP SELECTION
            convlowerable = IsConvLowerable(inst);   // sub_14553620
            conv          = ExtractConvLikeHlo(inst);// sub_1D6AA140
            priced        = inst;
            if conv and conv.opcode == 94 (reduce-window):
                t = GetReduceWindowType(conv);       // sub_1454D4A0
                if t != -1 and t != 2:               // 0/1 (lane/sublane) → conv-priced
                    if convlowerable:
                        return GetOutputFusionOrConvolutionCycles(inst, opts, fs);  // sub_130AEDE0
                    priced = inst;                   // fall to collective-compute / default
            else if convlowerable:
                return GetOutputFusionOrConvolutionCycles(inst, opts, fs);

            // (4) COLLECTIVE-COMPUTE FUSION (collective fused with compute)
            if IsCollectiveComputeFusion(priced):    // sub_13E028C0
                return GetCollectiveComputeFusionCycles(priced, opts, fs);  // sub_130B13A0
            // (5) DEFAULT
            return GetLoopFusionOrUnfusedHloCycles(priced, opts, fs);       // sub_130B2BC0

    // non-numeric / unsupported types → empty ResourceVector (Ok, all-zero)
    return Ok(empty_rv);

塑造分派的三个门

numeric element-type 门（@0x130aa9dc）。 在任何 conv/fusion 工作之前，分派先用 cmp et, 0x21 后接 bt 0x2FFF91FFE（@0x130aa9e2 的 movabs）测试 inst.shape.element_type。0x2FFF91FFE 是覆盖 PrimitiveType 序号的 34-bit mask，选择 bundle 模型可以定价的 numeric/packed 类型；tuple/token/opaque 类型（mask 外）会直接落到空（全零）ResourceVector。两个次级测试放宽该门：(et & 0xFFFFFFFE) == 0x20 接纳 F8 对（序号 0x20/0x21），而在 et ≤ 0x22 上的 bt 0x400048000 是 loop-fusion / output-fusion 剥离 mask，即使 fusion root 自身元素类型是结构性的，也允许它通过。

陷阱 — 同一个 0x2FFF91FFE numeric mask 和同一个 0x400048000 剥离 mask 会在 RecordHloCycles（@0x130bc03d）中逐字重现。它们不是同一个决策：GetHloResourcesImpl 中的副本门控路由（该 op 是否会到达 conv/fusion 分支），而 RecordHloCycles 中的副本门控 leaf 存入（该 op 是否被记录或短路）。重新实现必须在两层都应用该 mask；只在顶层应用会让结构性 op 滑入 leaf switch，只在 leaf 应用则会在顶层误路由 numeric op。

conv/reduce-window 选择（@0x130aaa79）。 ExtractConvLikeHlo 从（可能 fused 的）指令中拉出 conv-like root。如果该 root 是 reduce-window（opcode 94 = 0x5e），GetReduceWindowType 会分类其轴。反编译在 t != -1 && t != 2 上分支：类型为 -1（not conv-lowerable）或 2（major-axis / max-pool）的 reduce-window 不会按 convolution 定价，而是落到默认 loop 路径。只有 lane-axis（0）或 sublane-axis（1）且同时 IsConvLowerable 的 reduce-window 会到达 GetOutputFusionOrConvolutionCycles。

注意 — 这里的类型 -1/2 sentinel 是跳过 MXU 定价的分派过滤器；它不是 leaf reduce-window 发射器。当一个 reduce-window 确实按 conv 定价时，它随后会到达 RecordReduceWindowCycles，在那里全部三种轴类型（lane/sublane/major）都会存入。上游 skip 逻辑不能带入 leaf。完整 pooling 成本在 Reduce-Window / Pooling Cost，它共享的 conv 状态在 ConvolutionCostState。

collective 谓词及其 fusion 剥离（@0x130aa580，第一个 IsSupportedCollectiveHlo 之后的行）。 裸 collective 会立即路由到 GetCollectiveCycles。当 op 是 fusion（opcode 61 = 0x3d）时，分派遍历 fused computation 的指令列表，寻找内部的 collective HLO（对每个 fused instr 调用 IsSupportedCollectiveHlo），如果找到且该 op 不是 collective-compute fusion，就通过 GetCollectiveCycles 为该内部 collective 定价。这是唯一会把被定价 op 从顶层指令重新指向内部指令的分支。

函数映射

Function	Address	Role
CostModel::GetHloResources	0x130aa560	public wrapper
CostModel::GetHloResourcesImpl	0x130aa580	5 路路由分派
IsSupportedCollectiveHlo	0x130aeda0	collective 谓词（分支 1）
GetCollectiveCycles	0x130abfc0	network-model 存入（rv[+8]）
IsFusionSupportedHlo	0x130abee0	numeric + fusion-support 门
IsConvLowerable	0x14553620	conv-lowerable 谓词
ExtractConvLikeHlo	0x1d6aa140	拉出 conv/reduce-window root
GetReduceWindowType	0x1454d4a0	−1/2 max-pool sentinel（轴类别）
GetOutputFusionOrConvolutionCycles	0x130aede0	分支 2 — MXU output-fusion model
IsCollectiveComputeFusion	0x13e028c0	分支 3 谓词
GetCollectiveComputeFusionCycles	0x130b13a0	分支 3
GetLoopFusionOrUnfusedHloCycles	0x130b2bc0	默认 → 阶段 2

---

阶段 2 — RecordCyclesIfFused，Fusion-Peel 路由器

目的

GetLoopFusionOrUnfusedHloCycles 组合 software-pipelined prologue/steady/tail bundle 成本（见 Bundle-Aware Cost），并调用按 op 存入机制。对于 fused region，该机制是 RecordCyclesIfFused：一个识别 fusion root 种类并分派到专用重型发射器的路由器，在任何 leaf op 被定价之前剥离结构性 fusion。未 fused 的 op 直接到达 leaf RecordHloCycles。

入口点

RecordCyclesIfFused (@0x130cc720)            ── fusion-root kind router
  ├─ IsLoopFusion?    → RecordLoopFusionCycles      @0x130b89a0  (recurse over fused leaves)
  ├─ IsConvLowerable? ── builds shared ConvCostState, then:
  │     ├─ IsOutputFusion? → RecordOutputFusionCycles  @0x130b86c0  (MXU output-fusion)
  │     ├─ opcode == 0x2b?  → RecordConvolutionCycles   @0x130ca6c0  (conv → Matmul/Matpush/Xlu)
  │     └─ (else, CHECK 0x5e) → RecordReduceWindowCycles @0x130c94e0 (pooling → VectorLoad/AluAny/Xlu)
  └─ (not conv-lowerable) → RecordHloCycles         @0x130bbfe0  (leaf per-op deposit)
```text

### 算法

```c
Status RecordCyclesIfFused(inst, fs, window, rv, isFused, nesting):  // sub_130CC720
    if IsLoopFusion(inst):                              // line 188
        return RecordLoopFusionCycles(inst, window, rv, …, nesting+1);  // sub_130B89A0
    if IsConvLowerable(inst):                           // line 256 — gates the window arms
        convCostState = GetConvolutionCostState(inst);  // + CalculateNestedConvolutionWindows
        if IsOutputFusion(inst):                        // line 434
            return RecordOutputFusionCycles(inst, &convCostState, …, nesting+1);  // sub_130B86C0
        if inst.opcode == 0x2b (convolution):           // line 513
            return RecordConvolutionCycles(inst, &convCostState, rv, nesting+1);  // sub_130CA6C0
        // else: CHECK(opcode == kReduceWindow) @line 526 (cost_model.cc:6236)
        return RecordReduceWindowCycles(inst, …);                      // sub_130C94E0 — line 536
    // not conv-lowerable: drop internal producer edges, then the leaf deposit
    return RecordHloCycles(inst, window, rv, fs, nesting+1);           // sub_130BBFE0 — line 706

剥离顺序是固定的，并由五个 call site 确认。首先测试 IsLoopFusion（它递归遍历其构成 leaf）。剩余重型分支（output-fusion、然后 conv（0x2b）、然后 reduce-window（0x5e））嵌套在单个 IsConvLowerable(inst) 门（line 256）内：当该门成立时，路由器物化一个由三个 window 分支共享的 ConvCostState（通过 GetConvolutionCostState / CalculateNestedConvolutionWindows），并按上述顺序分派。只有非 conv-lowerable、非 loop-fusion op 才到达 leaf fallthrough。reduce-window 分支是 conv 测试的 else，并用 CHECK(hlo_to_fuse->opcode() == kReduceWindow) 自卫（MakeCheckOpString @line 526，cost_model.cc:6236），因此误分类剥离会 abort，而不是误定价。

注意 — leaf RecordHloCycles 自身会 CHECK !hlo->IsLoopFusion() && !hlo->IsOutputFusion()（cost_model.cc:6609，@0x130bbfe0 line 128）。fusion root 绝不能到达 leaf；先剥离它是 RecordCyclesIfFused 的职责。对 fusion root 调用 leaf 的重新实现会触发该断言。剥离路由器是从 fusion 到 deposit 的唯一路径。

边内部化 — IsProducerUse

在 fused parameter 的 leaf 存入之前，RecordHloCycles 会遍历 fusion 的 operands，并逐 operand edge 咨询 IsProducerUse（@0x130ab0c0）。fusion 内部的 producer→consumer 边会被丢弃 input-DMA；这些字节从不离开 fusion，因此不产生成本为 MemXfer 的周期。这是 bundle 侧对 fusion-priority 模型中边内部化的镜像：只有 fusion 的外部输入按 input transfer 收费。该遍历发生在 leaf switch 的 parameter 分支（下文），不在 RecordCyclesIfFused 中。

函数映射

Function	Address	Role
CostModel::RecordCyclesIfFused	0x130cc720	fusion-root kind router
CostModel::RecordLoopFusionCycles	0x130b89a0	loop-fusion recursion
CostModel::RecordOutputFusionCycles	0x130b86c0	output-fusion（MXU）
CostModel::RecordConvolutionCycles	0x130ca6c0	conv peel → MXU deposit
CostModel::RecordReduceWindowCycles	0x130c94e0	reduce-window peel → vector deposit
CostModel::RecordHloCycles	0x130bbfe0	leaf per-op deposit
CostModel::IsProducerUse	0x130ab0c0	producer→consumer edge-drop
CostModel::RecordFusionInputCycles	0x130ce940	external-input DMA → MemXfer

---

阶段 3 — RecordHloCycles，Leaf Opcode→Slot 路由

目的

RecordHloCycles 是 leaf 按 op 存入：给定一条 HLO 指令、bundle Window、输出 ResourceVector 和 nesting depth，它把该 op 的吞吐周期存入该 op 占用的具名 Resource 槽。op→block 决策是对 opcode byte（*(byte*)(a2+12)）的 switch，编译为 .rodata 0xae0ebbc 处 0x7f 项自相对 jump table。分派索引是 idx = opcode − 3（@0x130bc0db 的 add eax, 0xFFFFFFFD），ja > 0x7e 路由到 default block。

入口点

RecordHloCycles (@0x130bbfe0)
  ├─ CHECK(!IsLoopFusion && !IsOutputFusion)      cost_model.cc:6609
  ├─ numeric-mask + IsLoopFusion/IsOutputFusion peel pre-check   @0x130bc03d
  ├─ ElementIsFloating(shape)                     ── the add/subtract lane gate
  ├─ element_count = Product(output dims)         ── xla::Product → [rbp-0x30]
  ├─ RecordHloCyclesIfTopLevel  @0x130cdd80        ── == 1 guard, else skip the switch
  └─ switch(opcode) → jump table @ .rodata 0xae0ebbc
```text

### 存入机制

每个具名 block 都解析一个槽和一个按代际吞吐量，然后累加：

```c
// per CT::Instruction k issued by the block:
slot = CycleTable::GetResource(k);                  // sub_1C89CE20 = dword_B438AEC[k]
thru = cycletable.vtable[+0x10](k);                 // GetCyclesForThroughput(k), per-gen
ResourceVector::Acc(slot, element_count × thru × W);// sub_1C89ADC0, W = fixed FP multiplier

element_count 是 Product(output dims)。GetResource 是一个单一且代际不变的扁平查找：dword_B438AEC[k]（4 B/entry，.rodata 0xb438aec），下文重新解码，并且与 Resource Enum 上的表相同。GetCyclesForThroughput 是按代际整数（Per-Opcode Cycle Constants）。Acc（@0x1c89adc0）用会 trap 的 ud1 检查 slot < 0x17（23），并执行 vector[slot] += cycles。

怪癖 — 对于 float add/subtract，block 通过 GetResource(CT) 解析槽（→ VectorAlu1）；但对于 integer add/subtract，槽被硬编码为 5（VectorAluAny），而吞吐量仍然是 thru(CT 0x12) / thru(CT 0x13)。槽和吞吐量是解耦的：float 门只重定向目标 lane，不重定向周期计数。按类型门控吞吐量（而不是 lane）的重新实现会错误定价整数算术。

算法

RecordHloCycles(inst, window, rv, fs, nesting):       // sub_130BBFE0
    CHECK(!inst.IsLoopFusion() && !inst.IsOutputFusion());     // line 128
    isFloat = ShapeUtil::ElementIsFloating(inst.shape);        // line 139
    elems   = Product(output_dims);                            // line 145 → [rbp-0x30]
    if RecordHloCyclesIfTopLevel(...) != 1: return status;     // line 147 (fusion-root guard)

    switch (*(byte*)(inst + 12)):                     // jump table @ .rodata 0xae0ebbc
      case 0x03 add:        slot = isFloat ? GetResource(CT 0x12) : 5;   // R4 / R5
                            Acc(slot, elems × thru(CT 0x12));
      case 0x7b subtract:   slot = isFloat ? GetResource(CT 0x13) : 5;   // R4 / R5
                            Acc(slot, elems × thru(CT 0x13));
      case 0x4b multiply:   Acc(GetResource(CT 0x14) /*R3*/, elems × thru(CT 0x14));
      case 0x32 divide:     /* 4-deposit reciprocal+mul micro-sequence */
            Acc(6 /*R6 VectorEup*/,        elems × thru(CT 0x18));
            Acc(GetResource(CT 0x14) /*R3*/, (elems × thru(CT 0x14)) × 3.0);   // 0xa2df930
            Acc(GetResource(CT 0x12) /*R4*/, (elems×2) × thru(CT 0x12));
            Acc(5 /*R5*/,                  elems × 9.0);                       // 0xa2deb40
      case 0x47 logistic:   /* 5-deposit sigmoid micro-sequence */
            Acc(GetResource(CT 0x12) /*R4*/, elems × thru(CT 0x12));
            Acc(GetResource(CT 0x14) /*…*/, (elems×2) × thru(CT 0x14));
            Acc(5 /*R5*/, elems);
            Acc(6 /*R6 VectorEup*/, … via CT 0x1a);
      case 0x38 erf:        if (inst.vtable[+312](11)):          // ext-precision predicate
                                Acc(6 /*R6 VectorEup*/, elems × thru(CT 0x11));  // single deposit
                            else: /* 4-deposit polynomial */
                                Acc(6 /*R6*/,                  elems × thru(CT 0x18));
                                Acc(GetResource(CT 0x14) /*R3*/, (elems × thru(CT 0x14)) × 16.0); // 0xa2df040
                                Acc(GetResource(CT 0x12) /*R4*/, (elems×2) × thru(CT 0x12));
                                Acc(5 /*R5*/,                  elems × 4.0);     // 0xa2de830
      case 0x2a convert:    if ElementHasBitWidth(inst.shape, 1):  // 1-bit PRED
                                Acc(5, elems × 2.0);               // packed-bool repack
                            else: /* fall to ZERO — numeric width-conversion free */
      case 0x6c select:     Acc(5, elems × 2.0);                   // read both branches + pred
      case 0x52 parameter:  if !IsFused(inst): /* not deposited here */
                            else: walk operands, IsProducerUse edge-drop,
                                  RecordCyclesIfFused per producer, RecordFusionInputCycles → MemXfer;
      case 0x5b reduce:     if IsFused(inst): Acc(5, elems);
                            else: w = GetInputWindow(inst); Acc(5, Product(w));  // operand window
      case 0x18,0x1a,0x27,0x29,0x43,0x61,0x81: /* ZERO — no deposit */
      default:              Acc(5 /*R5 VectorAluAny*/, elems × 1.0);
```text

### 解码后的 Jump Table

Block 地址和 jump-table target 偏移已从 `.rodata 0xae0ebbc` 逐字节验证（`target = 0xae0ebbc + offset[opcode−3]`）。`divide`/`logistic`/`erf` 的完整 µop-sequence 细节（FP multiplier 缩放、slow-vs-fast erf 门）记录在 [TpuHloCostAnalysis](tpu-hlo-cost-analysis.md#recordhlocycles--the-opcodeslot-jump-table)；本表是 routing-spine 视图，即哪个 block、哪个 CT、哪个 slot。

| Opcode（name） | Block @ | 发出的 CT → slot | 周期量 |
|---|---|---|---|
| `0x03` add | `0x130bc0f7` | float: `CT 0x12`→`R4`; int: `R5` | `elems × thru(0x12)` |
| `0x7b` subtract | `0x130bc4af` | float: `CT 0x13`→`R4`; int: `R5` | `elems × thru(0x13)` |
| `0x4b` multiply | `0x130bc4dd` | `CT 0x14`→`R3` | `elems × thru(0x14)` |
| `0x32` divide | `0x130bc3cf` | `R6` + `CT 0x14`→`R3` + `CT 0x12`→`R4` + `R5` | 4-deposit；`×3.0`(R3)，`×2`(R4 elems)，`×9.0`(R5) |
| `0x47` logistic | `0x130bc17b` | `CT 0x12`→`R4`, `CT 0x14`, `R5`, `CT 0x1a`→`R6` | 5-deposit sigmoid；`×2`(R-elems) |
| `0x38` erf | `0x130bc245` | fast: `CT 0x11`→`R6`; slow: `R6`+`R3`+`R4`+`R5` | gated；slow `×16.0`(R3)，`×2`(R4)，`×4.0`(R5) |
| `0x2a` convert | `0x130bc293` | 1-bit: `R5`; wider: none | 1-bit: `elems × 2.0`；wider: 0 |
| `0x6c` select | `0x130bc2a9` | `R5`（`×2` path） | `elems × 2.0` |
| `0x52` parameter | `0x130bc2b8` | fused→`RecordFusionInputCycles`; else none | fused: input-DMA into `MemXfer` `R9..12` |
| `0x5b` reduce | `0x130bc2f9` | `R5` over operand window | 按 reduced-OVER input window 定价 |
| DEFAULT（大多数 op） | `0x130bc3c0` | `R5` `VectorAluAny` | `elems × 1.0` |
| ZERO-cost | `0x130bc8c2` |（无存入）| `0` |

DEFAULT 分支覆盖上面未具名的每个 numeric elementwise / structural op（反编译器注释为 `cases 4-23, 25, 27-38, 40, 43-49, 51-55, 57-66, 68-70, 72-74, 76-81, 83-90, 92-96, 98-107, 109-122, 124-128`）。ZERO 分支覆盖 data-layout op：`bitcast`（`0x18`）、`broadcast`（`0x1a`）、`concatenate`（`0x27`）、`constant`（`0x29`）、`iota`（`0x43`）、`reshape`（`0x61`）、`tuple`（`0x81`），全部位于单个 block `0x130bc8c2`。

> **注意 —** `erf` fast path 不携带 `Matpush` traffic：ext-precision 谓词 `inst.vtable[+312](11)`（`@0x130bc245`）选择单个 `VectorEup` 存入 `CT 0x11`，`GetResource` 将其映射到 `R[6]`（`dword_B438AEC[0x11] == 6`，逐字节验证）。它完全停留在 vector/EUP pipe 上；见 [TpuHloCostAnalysis](tpu-hlo-cost-analysis.md)。

### Op→Slot 表 — `dword_B438AEC`

`CycleTable::GetResource(k)`（`@0x1c89ce20`）字面上就是 `return dword_B438AEC[k]`。leaf block 发出的条目，从二进制重新解码如下：

| `CT::Instruction` | `dword_B438AEC[CT]` → slot | 发出者 |
|---|---|---|
| `0x11` | `R[6]` `VectorEup` | erf fast path |
| `0x12` | `R[4]` `VectorAlu1` | float add, divide, logistic |
| `0x13` | `R[4]` `VectorAlu1` | float subtract |
| `0x14` | `R[3]` `VectorAlu0` | multiply, divide, logistic, erf-slow |
| `0x18` | `R[6]` `VectorEup` | divide / erf-slow transpose stage |
| `0x1a` | `R[6]` `VectorEup` | logistic lane-compare |

> **怪癖 —** `CT 0x12` 和 `CT 0x13` 都映射到*同一个*槽 `R[4]` `VectorAlu1`，尽管它们是不同的 `CycleTable::Instruction` 序号（input-rotate 与 output-rotate），并且有不同的按代际吞吐量。因此 float add 和 float subtract 共享专用 lane，但周期计数可能不同。Multiply 固定到*另一个*专用 lane（通过 `CT 0x14` 到 `R[3]`）。这是 `MaxResourceCycles` 的有意端口均衡输入；见注意事项。

### 函数映射

| Function | Address | Role |
|---|---|---|
| `CostModel::RecordHloCycles` | `0x130bbfe0` | leaf per-op deposit + opcode switch |
| `CostModel::RecordHloCyclesIfTopLevel` | `0x130cdd80` | `== 1` fusion-root guard |
| `CycleTable::GetResource` | `0x1c89ce20` | `dword_B438AEC[CT]`（代际不变） |
| `ResourceVector::Acc` | `0x1c89adc0` | `vector[slot] += cycles`（边界 `< 23`） |
| `ShapeUtil::ElementIsFloating` | (call @line 139) | add/subtract lane gate |
| `ShapeUtil::ElementHasBitWidth` | `0x20ce1580` | 1-bit PRED convert gate |
| `xla::Product` | `0x20cf5200` | `element_count` = Product(output dims) |
| `GetInputWindow` | (anon-ns) | non-fused reduce operand window |

### 固定 FP Multiplier（`.rodata`，逐字节验证）

| Address | Value | Used by |
|---|---|---|
| `0xa2df230` | `1.0` | default per-op multiplier |
| `0xa2df930` | `3.0` | divide `VectorAlu0`（R3）scale |
| `0xa2deb40` | `9.0` | divide `VectorAluAny`（R5）scale |
| `0xa2df040` | `16.0` | erf slow-path `VectorAlu0`（R3）scale |
| `0xa2de830` | `4.0` | erf slow-path `VectorAluAny`（R5）scale |
| `0xa2df5c8` | `0.5` | `MaxResourceCycles` `VectorAlu` port-balance blend |

---

## 注意事项 — 路由如何输入 Reduction

leaf op→slot 路由跨三个 `VectorAlu` 槽分布，是 `MaxResourceCycles` 0.5-blend group `{R3, R4, R5}` 的端口均衡输入（见 [Resource Enum](resource-enum.md#the-maxresourcecycles-reduction)）。该路由是有意的：multiply 固定到 `VectorAlu0`（`R3`），floating add/subtract 固定到 `VectorAlu1`（`R4`），而 catch-all DEFAULT、`×2` op（`select`、1-bit `convert`）以及 integer add/subtract 落在 `VectorAluAny`（`R5`），即 blend 会重新分配的 lane。因此混合 multiplies、floating adds 和 “any” ops 的 fusion 会填满两个专用 lane，bundle 成本是 balanced max，而不是串行总和。

阶段 1 和 2 到达的重型发射器会存入*其他*槽组。`RecordConvolutionCycles`（conv peel）存入 `Matmul` / `Matpush` / `Xlu`（`R0`/`R1`/`R2`，plain-MAX MXU group，见 [ConvolutionCostState](convolution-cost-state.md)）；`RecordReduceWindowCycles`（pooling）存入 `VectorLoad` / `VectorAluAny` / `Xlu`（`R7`/`R5`/`R2`，见 [Reduce-Window / Pooling Cost](reduce-window-pooling-cost.md)）；`GetCollectiveCycles` 为 ICI links 将一个标量 cycle count 写入 `rv[+8]`。因此每个路由分支都会落在 reduction rule 与该 op 硬件行为匹配的槽组：MXU pipe 重叠（plain MAX），vector ALU 端口均衡（0.5 blend），memory transfer 串行化（sum）。

`parameter` 和 `reduce` 会从 ALU lane *路由出去*。fused `parameter` 是 external-input DMA，由 `RecordFusionInputCycles`（`@0x130ce940`）在 `IsProducerUse` 丢弃内部 producer edge 之后定价到 `MemXfer` 槽 `R9..12`；VMEM-resident 参数是否改按 memory space 路由到 `R7` `VectorLoad` 尚未逐字节确认（MEDIUM）。non-fused `reduce` 通过 `GetInputWindow` 按其*operand* window 定价，成本随大的 reduced-over tensor 缩放，而不是随小的输出缩放；这是 bundle 侧对 `HandleReduce` 的 `ExtentProduct(operand)` flop 公式的镜像。

> **怪癖 —** bundle 模型在 leaf 中把 `broadcast`（`0x1a`）定价为**零**（ZERO-cost block），即使 fusion-priority 模型可能为同一个 op 收取 cross-lane-movement weight。这两个表面对 layout op 的分歧是有意设计：bundle 路径把 broadcast 视为免费的寄存器 splat（无功能单元占用），而 priority 模型会计入数据移动。不要统一它们。见 [NormalizedComputationCost](normalized-computation-cost.md)。

---

## 示例 — 一个 Fusion 通过全部三个阶段的路由

一个 loop-fusion body `{ multiply [256,128] (f32), add [256,128] (f32), tanh [256,128] }` 自顶向下定价：

```text
GetHloResourcesImpl(fusion)                         Stage 1
  IsSupportedCollectiveHlo → no
  et = F32 (in 0x2FFF91FFE mask) → numeric gate OK
  IsFusionSupportedHlo → yes; ExtractConvLikeHlo → null (no conv)
  IsConvLowerable → no; IsCollectiveComputeFusion → no
  → GetLoopFusionOrUnfusedHloCycles                 (default arm)

GetLoopFusionOrUnfusedHloCycles → RecordCyclesIfFused(fusion)   Stage 2
  IsLoopFusion → yes → RecordLoopFusionCycles → recurse over leaves:
    each leaf → RecordCyclesIfFused → (not a fusion/conv/rw) → RecordHloCycles

RecordHloCycles per leaf, element_count = 256·128 = 32768       Stage 3
  multiply (0x4b)  → CT 0x14 → GetResource → R3 VectorAlu0:  Acc(R3, 32768 × thru(0x14))
  add (0x03), f32  → CT 0x12 → GetResource → R4 VectorAlu1:  Acc(R4, 32768 × thru(0x12))
  tanh (0x7d)      → DEFAULT block →           R5 VectorAluAny: Acc(R5, 32768 × 1.0)

MaxResourceCycles 在 {R3, R4, R5} 上的 0.5-blend 会让 multiply（VectorAlu0）和 add（VectorAlu1）lane 并行运行，把 tanh 的 “any” 工作负载均衡到较空闲的 lane，并以 50% 重叠剩余部分；所以 bundle 成本是 ≈ max(R3, R4) + half the tanh residual，而不是串行总和。如果 add 是 integer 类型，它会与 tanh 一起进入 R5（当 ElementIsFloating 为 false 时硬编码为槽 5），而不是固定到 R4。flop 侧（由单独 pass 运行）将 tanh 分类为 transcendental（Properties+0x54），将 multiply/add 分类为 flops（Properties+0x50）；两个表面对哪些 op 昂贵达成一致，但从不共享数值。

---

相关组件

Component	Relationship
TpuHloCostAnalysis	与此路由并行运行的 flop/byte override 表面；拥有完整的 RecordHloCycles µop-sequence 细节和 flop-override 表
Resource Enum (23-slot)	ResourceVector 槽、Acc、GetResource，以及 deposits 输入的 MaxResourceCycles reduction
ConvolutionCostState	从阶段 1/2 到达的 conv peel 发射器（RecordConvolutionCycles）
Reduce-Window / Pooling Cost	reduce-window peel 发射器和 GetReduceWindowType 轴分类器
Bundle-Aware Cost	覆盖按 op deposits 的 software-pipelined loop 组合
NormalizedComputationCost	对同一 HLO 做不同定价的 fusion-priority 表面

交叉引用

• TpuHloCostAnalysis — 输入此路由主干的 override 表面；flop/byte Properties 模型以及按 op µop-sequence 分解
• Resource Enum (23-slot) — Resource 槽名、Acc、GetResource op→slot 表，以及 MaxResourceCycles overlap 模型
• ConvolutionCostState — 从 conv 分支到达的 RecordConvolutionCycles / RecordConvKernelCycles，即 MXU deposit
• Reduce-Window / Pooling Cost — RecordReduceWindowCycles 和 lane/sublane/major 轴分派
• WindowDescription Byte-Cost — 重型发射器背后的 conv/DMA byte+throughput 原语
• NormalizedComputationCost — 对 layout op 定价不同于 bundle 路径的 opcode→weight fusion-priority 模型
• Bundle-Aware Cost — GetLoopFusionOrUnfusedHloCycles 以及覆盖按 op deposits 的 prologue/steady/tail loop 成本
• Per-Opcode Cycle Constants — 每个 deposit 乘以的按代际 GetCyclesForThroughput(CT::Instruction) 整数
• Cost Model Overview — 此路由在成本模型类族中的位置