Internal Wiki

主题

kDeviceTypeInfo 生产者与 Roofline 读取者 ​

本页所有地址适用于 libtpu-0.0.40-cp314 wheel 中的 libtpu.so（build-id 89edbbe81c5b328a958fe628a9f2207d，这是明确锚点；运行时报告的 0.103 无法在二进制中静态验证）。该二进制 未 strip；每个符号都是 demangled C++ 名称。.text、.rodata 和 .lrodata 映射为 VMA == 文件偏移；kDeviceTypeInfo 位于 .lrodata 中的 VMA 0x1c60480。其他 build 会不同。

摘要 ​

xprof::kDeviceTypeInfo 是 profiler 的按世代硬件规格表：一个 17 × 0x448 字节的 DeviceTypeInfo 记录数组，每个 DeviceType ordinal 一条，保存各精度峰值 FLOP/s、各 memory-space 带宽、时钟、core 几何、DVFS 阶梯、perf-counter-set descriptor，以及 firmware power/thermal bundle。姊妹页面 按 DeviceType 的 Profiler 结构 负责字节级字段布局；v7x Perf-Counters 负责 counter-name descriptor 字段。本页负责另外两个事实：谁写入这张表，以及谁读取 roofline 相关数值。

producer 故事短而确定：没有运行时 producer。该表是一个通过 lea 寻址的标量 aggregate，位于 .lrodata，这是映射为只读不可写（perm 4）的 segment；因此没有 .text store 能指向它，loader 也不会对它应用 relocation（[0x1c60480, 0x1c64d48) 中有零个 .rela.dyn 条目）。每个字段都是链接时冻结的内联 int32、int32[8] 或 IEEE-754 double literal。所谓 “producer” 是发出 static const DeviceTypeInfo[17] initializer 的 C++ 前端；散布在 .lrodata 中的 13 个逐字节相同副本，是按实例化引用产生的 identical-code-folding（ICF）数据副本，而不是独立的表。

reader 故事则是 roofline 所在处；关键发现是 libtpu 不计算 roofline。它 盖印 roofline 的输入。三类 reader 拉取每行字段，并把它们作为 device-capability XStat 写到 XPlane 上：ConvertTpuTraceToXPlaneV2 发射 peak-TFLOP/s 和按 memory-space 划分的 bandwidth block（对每个 bandwidth 成员执行 base-10 GB/s → GiB/s 转换）；ConvertTpuTraceToXPlane 将 v7x perf-counter-set mask 馈入 counter-name resolver；ConvertFirmwareTraceEntriesToXPlane 将 power/thermal coefficient bundle 馈入 FirmwareEventBuilder。第四个 reader，即 cost model 的 XProfTpuCostAnalysis::HandleConvolution，是二进制内唯一对 peak-FLOP/s 字段做 roofline 风格 算术 的消费者，即按 operand precision 取最小值再除以 core count。+0xb8 以下不在这些路径上的所有内容（effective-HBM-bandwidth、latency、secondary-rate double）完全没有二进制内 reader；host 侧 xprof roofline 工具会在它们被盖印后消费。

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

  • producer 模型：只读 .lrodata 中编译期 static const DeviceTypeInfo[17]，按 base + ordinal*0x448 索引，没有 writer、没有 relocation、没有 pointer 字段；通过 lea GOT; movabs (copy - GOT); add; imul ordinal,0x448 寻址。
  • device-capability stat-stamp 路径：StatType 0x62..0x6b → table-offset 映射、应用到五个 bandwidth stat 的 ×1.073741824 GB→GiB 因子，以及它们填充的 peak_* stat-string 集。
  • cost-model reader：HandleConvolution 对 +0x60/+0x78/+0x80 逐 element-type 执行 vminsd 并除以 core count，来源于 xprof_tpu_cost_analysis.cc:190。
  • producer→table→reader provenance：kDeviceTypeInfo 是一个 独立的 profiler spec table，不是 compiler 的 chip_parts FlopsPerSecond 的字节冻结副本；二者共享硬件谱系，但按世代会分歧。
表符号xprof::kDeviceTypeInfo (_ZN5xprofL15kDeviceTypeInfoE) @ 0x1c60480
布局17 × 0x448 = 18632 字节；LOCAL OBJECT；结束于 0x1c64d48
Section / perm.lrodata [0x1884a00, 0x84931d0)；perm 4（R，无 W，无 X）
Producer编译期 static const aggregate；无运行时 writer；零 relocation
ICF 副本13 个逐字节相同 data fold；按 base + ordinal*0x448 寻址
Cap-stat readerConvertTpuTraceToXPlaneV2<…jxc>::lambda @ 0xf1da7c0
Cost readerXProfTpuCostAnalysis::HandleConvolution @ 0xf58e7c0
Perf-set readerConvertTpuTraceToXPlane<…jxc> @ 0xf23f8c0
Firmware readerConvertFirmwareTraceEntriesToXPlane<…vlc> @ 0xf27f140
GB→GiB 因子qword_A2E0210 = 1.073741824 (= 2³⁰ / 10⁹)

Producer — 编译期 Static Const，无运行时 Writer ​

目的 ​

重新实现者的第一个问题是 kDeviceTypeInfo 是在进程启动时由 constructor 填充（那就必须复现），还是被烘进镜像（只需转录）。二进制给出答案：烘进镜像。libtpu 中没有 producer 函数；这些字节是链接时常量。

证据 ​

三个独立事实各自都足够；合在一起则是决定性的。

text
1. Section + permission.  kDeviceTypeInfo @0x1c60480 lies in .lrodata,
   mapped [0x1884a00, 0x84931d0) with perm 4 (has_write=False, has_code=False).
   A runtime store to 0x1c60480 would fault — the page is never writable.

2. Relocations.  .rela.dyn carries ZERO entries in [0x1c60480, 0x1c64d48).
   No field is a relocated pointer; every field is an inline scalar/float
   literal.  ⇒ no codename-string ptr, no per-gen sub-table ptr — nothing
   the loader patches, nothing a constructor fills.

3. Symbol shape.  nm reports a LOCAL OBJECT of size 18632 = 17 × 0x448,
   exactly the 17-entry stride.  A runtime-built table would be in .bss
   or .data (writable); this one is rodata-class.
```text

该符号在名称表中可清晰解析：

```text
_ZN5xprofL15kDeviceTypeInfoE  → xprof::kDeviceTypeInfo  (13 entries, all same name, size 0x48c8 each)
   0x1c3a910 0x1c3f1e0 0x1c43ab0 0x1c48380 0x1c4e140 0x1c52a10 0x1c572e0
   0x1c5bbb0 0x1c60480 0x1c7bbf0 0x1c804c0 0x1c84d90 0x1c89660

nm 报告十三个 LOCAL OBJECT 条目，全部携带 相同 的 mangled name _ZN5xprofL15kDeviceTypeInfoE（没有 _0/_1 disambiguator 后缀），且大小相同为 0x48c8；0x1c60480 是本页 reader VA 解析到的副本。mangling 中的 L（...L15kDeviceTypeInfoE）标记它为 static（internal linkage），也就是 translation-unit-local const array，正是 static const DeviceTypeInfo kDeviceTypeInfo[17] = { … } 的产物。因为它是 internal-linkage 并由许多不同的 template/pass 实例化引用，linker 的 ICF pass 将相同 data 折叠成副本，而不是把它们合并为单个符号；这 13 个副本是 folding artifact，不是 13 张不同的表。

注意 — “没有 producer 函数”并不意味着“没有 producer”。producer 是发出 static const aggregate initializer 的 compiler。重新实现应把它复现为 constexpr/static const 数组，而不是 __attribute__((constructor)) 或 lazy-init singleton。把它当作运行时初始化会增加一个原始实现不存在的 writer。

寻址模式 — Reader 如何找到一行 ​

每个 reader 以相同方式到达一行。表基址相对 GOT materialize（因此同一段代码可使用 linker 分配给该 translation unit 的任意 ICF 副本），并通过将 DeviceType ordinal 乘以 0x448 stride 来选择行。

c
// The canonical row-load, as it appears in every reader (e.g. 0xf1da7c0):
base = &GLOBAL_OFFSET_TABLE_ + (copy_VMA - GOT_anchor) / 8;   // lea + movabs fold
if (ordinal >= 0x11) trap();                                  // 17-entry bound check
row  = (char*)base + 0x448 * ordinal;                         // imul ordinal, 0x448
field = *(double*)(row + disp);                               // vmovsd [base+idx+disp]
```text

> **陷阱 —** IDA decompiler 会把 stride 渲染为基于 8 字节元素指针的 `137 * ordinal`（`137 × 8 = 0x448`）或基于 4 字节元素指针的 `274 * ordinal`（`274 × 4 = 0x448`）。二者都是同一个 `0x448` 字节 stride；不要把 `137`/`274` 解读为元素数量。反复出现的 `ordinal >= 0x11` guard（其上方有 `ud1` trap）是 17 行边界检查，存在于每个加载点。

### 来源 — 独立于 `chip_parts`，不是副本

compiler 的按世代硬件常量位于嵌入式 `chip_parts` proto 中（见 [chip_parts.binarypb](../targets/chip-parts-binarypb.md)）；profiler 的常量在这里。它们是 **只由 `DeviceType` ordinal 连接的两个平行 spec store**，而不是一个冻结副本。解码 `chip_parts` `FlopsPerSecond` 表并与 `kDeviceTypeInfo +0x60/+0x78/+0x80` peak double 比较，会看到 bf16 数值和 *所有* v5p 槽位分歧：

| 槽位（`kDeviceTypeInfo` 偏移） | `kDeviceTypeInfo` TFLOP/s | `chip_parts` /1e12 | 比率 | 匹配 |
|---|---|---|---|---|
| v5p bf16 `+0x60` | 236.7 | 197.0 | 1.202 | NO |
| v5p int8 `+0x78` | 466.2 | 394.0 | 1.183 | NO |
| v5p int4 `+0x80` | 925.2 | 788.0 | 1.174 | NO |
| v6e bf16 `+0x60` | 946.7 | 918.0 | 1.031 | NO |
| v6e int8 `+0x78` | 1835.0 | 1835.0 | 1.000 | YES |
| v6e int4 `+0x80` | 3670.0 | 3670.0 | 1.000 | YES |

只有 v6e int8/int4 槽位逐字节一致；这是因为二者都来自相同的 `systolic_dim² × clock` 硬件规格，而不是因为 profiler 复制了 compiler blob。bf16 数值跟随更高时钟推导（v6e `+0x60` 946.7 暗示原始 `256² × clock` 数值，而 `chip_parts` 918 是扣除开销后的公开数值）。对重新实现者来说，这意味着 profiler 的 roofline ceiling 与 compiler 的 cost model 可以按世代合法地不同；不能把它们合并。

---

## Reader 1 — Device-Capability Stat Block（`ConvertTpuTraceToXPlaneV2`）

### 目的

这是核心 roofline-input producer。profiler 将 TPU device trace 转换为 XPlane 时，会在 device plane 上盖印一组 device-capability `XStat`：peak TFLOP/s 和各 memory-space 的 peak bandwidth。host 侧 xprof roofline 工具从 XPlane 读回这些 stat 来绘制 roofline。libtpu 的工作止于盖印；它自己不做 ceiling-vs-bandwidth 分类。

### 入口点

```text
ConvertTpuTraceToXPlaneV2<…jxc::PerformanceTraceEntry>   ── device-trace → XPlane (F-copy)
  └─ lambda @0xf1da7c0                                    ── stamps the capability stat block
       ├─ GetStatTypeStr(98..108)                         ── 0x1cf8c420, StatType → metadata key
       ├─ vmovsd [base + idx*0x448 + disp]                ── load the per-row spec double
       ├─ vmulsd ×qword_A2E0210 (=1.073741824)            ── GB/s → GiB/s, bandwidth members only
       └─ XStatsBuilder::SetStatValue                     ── write the XStat onto the plane

算法 ​

该 lambda 遍历固定的一段 stat ID。对每个 ID，它通过 GetStatTypeStr(id) 获取 stat-name metadata，加载一个 kDeviceTypeInfo 字段，可选地乘以 GB→GiB 因子，然后追加 XStat。

c
function StampDeviceCapabilities(plane_builder, device_type):     // 0xf1da7c0
    bound_check(device_type < 0x11)                               // 17-row guard
    base = kDeviceTypeInfo                                        // GOT-relative ICF fold
    row  = base + device_type * 0x448

    emit(0x62, row[+0x60])                          // peak_teraflops_per_second  — no scaling
    emit(0x63, row[+0xd0] * GiB)                    // peak_hbm_bw_…              — ×1.073741824
    emit(0x64, row[+0x120] * GiB)                   // peak_*_bw (CMEM-class, v4) — ×GiB
    emit(0x65, row[+0x128] * GiB)                   // peak_*_bw (CMEM-class, v4) — ×GiB
    emit(0x66, row[+0x100] * GiB)                   // peak_sram/vmem_*_bw_…      — ×GiB
    emit(0x67, row[+0x108] * GiB)                   // peak_sram/vmem_*_bw_…      — ×GiB
    emit_byte(0x6b, byte(row[+0x2c4]))              // has_megacore               — BYTE flag
    // 0x68/0x69 read a DIFFERENT sub-indexed slot (see GOTCHA below), not this ladder
    // 0x6c.. continue with non-table stats (core counts, gtc, …)
```text

GB→GiB 因子 `1.073741824` 恰好是 `2³⁰ / 10⁹`：spec table 存储 base-10 GB/s（`10⁹` bytes/s），stat 以 GiB/s（每 GiB 为 `2³⁰` bytes）即 bytes/ns 发射。peak-TFLOP/s stat（`0x62`）*不*缩放，因为 TFLOP/s 已经是 base-10 速率，不涉及二进制前缀重解释。

### StatType → Offset 映射

每一行都由所引 site 的反汇编 `vmovsd` displacement 确认；当且仅当该成员是 bandwidth 时存在乘法。

| StatType | `GetStatTypeStr` 参数 | 表偏移 | 缩放 | Stat string（`.rodata`） | Site |
|---|---|---|---|---|---|
| `0x62` | 98 | `+0x60` | 无 | `peak_teraflops_per_second` (`0x86ed48f`) | `0xf1daa9b` |
| `0x63` | 99 | `+0xd0` | ×GiB | `peak_hbm_bw_gigabytes_per_second` (`0x86ed53b`) | `0xf1dab40` |
| `0x64` | 100 | `+0x120` | ×GiB | 一个 `peak_*_bw_…`（CMEM-class，仅 v4 非零） | `0xf1dabe6` |
| `0x65` | 101 | `+0x128` | ×GiB | 一个 `peak_*_bw_…`（CMEM-class，仅 v4 非零） | `0xf1dac8c` |
| `0x66` | 102 | `+0x100` | ×GiB | 一个 `peak_sram/vmem_*_bw_…` | `0xf1dad32` |
| `0x67` | 103 | `+0x108` | ×GiB | 一个 `peak_sram/vmem_*_bw_…` | `0xf1dadd8` |
| `0x6b` | 107 | `+0x2c4`（BYTE） | 无 | `has_megacore` (`0x86a4474`) | `0xf1db020` |

八个 `peak_*` stat string 在 `.rodata` 中连续排列：

```text
0x86ed48f  peak_teraflops_per_second
0x86ed4cc  peak_vmem_wr_bw_gigabytes_per_second
0x86ed4f1  peak_cmem_wr_bw_gigabytes_per_second
0x86ed516  peak_sram_wr_bw_gigabytes_per_second
0x86ed53b  peak_hbm_bw_gigabytes_per_second
0x86ed55c  peak_vmem_rd_bw_gigabytes_per_second
0x86ed581  peak_cmem_rd_bw_gigabytes_per_second
0x86ed5a6  peak_sram_rd_bw_gigabytes_per_second

+0xd0 HBM-bandwidth 成员是权威字段，并与 compiler 的 chip_parts HBM bandwidth 逐字节相符：v7x +0xd0 = 3433 GB/s × 1.073741824 = 3686.2 GiB/s == chip_parts 3.686 TB/s；v6e 1525.5 × 1.073741824 = 1638.0 GiB/s == 1.638 TB/s。这个 cross-check 将 +0xd0 锚定为 peak_hbm_bw_gigabytes_per_second，而不是 value-scaled +0xb8 组（这里 没有 读取该组）。

注意 — +0xf8–+0x130 组中被读取的成员（+0x100/+0x108/+0x120/+0x128）是 以 GB/s 为单位的 per-memory-space bandwidth，不是 peak-count 字段。每个成员在盖印前都乘以 1.073741824 GiB 因子；v4-only 的 +0x120/+0x128 与 CMEM bandwidth 对齐（v4 是唯一的 CMEM 世代）。+0xf8/+0x110/+0x118/+0x130 在此 build 中未被读取。

陷阱 — StatType 0x68/0x69（GetStatTypeStr(104)/(105)）不读取简单的 row + disp bandwidth ladder。它们计算 &base[0x448*idx] + 32*lut[idx-7] + {0xa0|0xa8}，其中 lut = qword_AB58328 是 7-entry 的按世代 sub-index。+0xa0/+0xa8 槽位在此 build 的整张表中为零，因此这两个 stat 发射零；若重新实现把 0x62..0x67 pattern 原样复制到 0x68/0x69，会错误寻址。

注意 — 0x64/0x65/0x66/0x67 分别绑定到六个 peak_{sram,vmem,cmem}_{rd,wr}_bw_… string 中哪一个，尚 未 在字节层面解析：GetStatTypeStr 是一个 lazy hash map，由 .text.startup 中的内联 initializer_list seed，而该 seed 顺序尚未解码。HBM（+0xd0）和 CMEM（+0x120/+0x128，仅 v4）绑定由数值 cross-check 明确；0x66/0x67 内部的 SRAM-vs-VMEM 拆分是中等置信度。

Reader 2 — Cost-Model Peak FLOP/s（HandleConvolution） ​

目的 ​

这是二进制内唯一对 peak-FLOP/s 字段执行 roofline 风格 算术 的 reader，而不只是盖印它们。XProfTpuCostAnalysis::HandleConvolution 通过对所有 operand precision 取硬件在该 precision 下提供的 最慢（最小）peak-FLOP/s，再除以 core count，来计算 convolution 的 peak achievable rate；这是 cost model 对该 op 计费的 per-op compute ceiling。

入口点 ​

text
XProfTpuCostAnalysis::HandleConvolution  @0xf58e7c0   ── src: xprof_tpu_cost_analysis.cc:190
  ├─ load row[+0x60]  (init, bf16/fp32-class ceiling)
  ├─ for each operand: switch(Shape::element_type - 2)
  │     ├─ int8-class  → vminsd against row[+0x78]   (≈2× +0x60)
  │     └─ int4-class  → vminsd against row[+0x80]   (≈4× +0x60)
  └─ vdivsd by core count / fixed divisors (6.0, 3.0)
```text

### 算法

```c
function HandleConvolution(this, conv):                       // 0xf58e7c0
    if not ShapeUtil::IsConvolution(conv):                    // 190 → AddSourceLocationImpl
        return error("xprof_tpu_cost_analysis.cc:190")
    idx   = this.device_type                                  // *((int*)this + 115)
    base  = kDeviceTypeInfo + idx*0x448
    peak  = base[+0x60]                                       // bf16/fp32-class peak (TFLOP/s)
    acc   = DBL_MAX                                           // qword_A2DEAF0 — vminsd seed
    for operand in conv.operands:
        et = Shape::element_type(operand) - 2                 // 28-way switch
        switch et:
            case int8-class:  acc = min(acc, base[+0x78])     // +0x78  (cmp ecx,6 sub-index)
            case int4-class:  acc = min(acc, base[+0x80])     // +0x80  (cmp ecx,3 sub-index)
            // fp32 / bf16 paths keep base[+0x60]
    rate = min(acc, peak)
    return rate / 6.0 / 3.0                                   // qword_A2DE720, qword_A2DF930 — ÷cores

vminsd-accumulate-over-operands 是关键思想：混合 bf16 input 与 int8 weight 的 convolution，只有当 int8 是 更慢 路径时才按 int8 ceiling 计费；running minimum 会选出硬件在各 precision 中维持最差的那个。÷6.0 ÷3.0 链是 per-core normalization（v7x 的 cores × logical-devices divisor）。+0x78 ≈ 2× 和 +0x80 ≈ 4× 相对于 +0x60 的逐字节比率（v5p/v5e/v6e 上分别为 1.94–1.97× 和 3.88–3.92×）确认它们是 int8/fp8 与 int4/fp4 peak-FLOP/s 槽位。

注意 — v7x 行在 +0x68 和 +0x80 都 存储 int8 数值 1992，而 +0x78 保存 996；v7x element-type dispatch 与 v5/v6 路径略有不同（它会把 bf16 operand 与 literal 1992.0、qword_A2DECF8 取 min）。+0x68 是备用 int8 槽还是 double-pumped fp8 rate 尚未解决（低置信度）；对 v5/v6，+0x78 是 2× 槽。

该 reader 与 XLA cost model 共享；周边分析见 TpuHloCostAnalysis，bandwidth 侧 cost 路径见 Memory Bandwidth & Latency Model。

Reader 3 — Perf-Counter-Set Mask（ConvertTpuTraceToXPlane） ​

目的 ​

G-copy converter 读取四个 +0x2c8/+0x348/+0x350/+0x358 dword（仅 v7x 行非零），并将每个作为 per-set enum base 传给 GetPerformanceCounterNames<N>；后者把用户选择的 counter ordinal 转换为盖印到 v7x counter 时间线上的片上寄存器名。本页只确立 表字段馈入 resolver；resolver 自身及其 DeviceType == 12 gate 由 v7x Perf-Counters 记录。

入口点 ​

text
ConvertTpuTraceToXPlane<…jxc::PerformanceTraceEntry>  @0xf23f8c0   (G-copy)
  ├─ GetPerformanceCounterNames<28>(dt, idxs, n>>1, row[+0x2c8], out)   @0xf240980
  ├─ GetPerformanceCounterNames<28>(dt, …,         row[+0x348], out)
  ├─ GetPerformanceCounterNames<28>(dt, …,         row[+0x350], out)
  ├─ GetPerformanceCounterNames<28>(dt, …,         row[+0x358], out)
  ├─ GetPerformanceCounterNames<3> (dt, …,         row[+0x440], out)    @0xf240ac0
  └─ GetPerformanceCounterNames<12>(dt, …,         row[+0x438], out)    @0xf240c00
```text

这六个字段作为 resolver 的 `base` 参数（`a4`）传入，resolver 用它计算 `PerformanceCounterNameToString(base + 8*ordinal)`。将六个 call-site GOT displacement 通过 `base + ordinal*0x448` 映射回来，正好解析到 `+0x2c8/+0x348/+0x350/+0x358`（四个 `<28>` TensorCore/SparseCore set）、`+0x440`（`<3>` CMNUR/HBM set）和 `+0x438`（`<12>` ICR set），也就是姊妹页面从 DT12 行 `0x1c637e0` 恢复出的同六个 descriptor 字段。

| Counter set | Resolver | 表字段 |
|---|---|---|
| TCS / SCS / SCTC / SCTD | `<28>` @ `0xf240980` | `+0x2c8` / `+0x348` / `+0x350` / `+0x358` |
| CMNUR (HBM/mem-net) | `<3>` @ `0xf240ac0` | `+0x440` |
| ICR (router) | `<12>` @ `0xf240c00` | `+0x438` |

---

## Reader 4 — Firmware Power/Thermal Bundle（`ConvertFirmwareTraceEntriesToXPlane`）

### 目的

firmware-trace converter 读取 `+0x360..+0x398` calibration bundle（四个 double 和四个 ulong），并将它们原样传给 `FirmwareEventBuilder` constructor；后者保存它们，并用它们在 firmware event 上盖印 `power` 和 `temperature` `XStat`（按世代的 DVFS/thermal model）。

### 入口点

```text
ConvertFirmwareTraceEntriesToXPlane<…vlc>  @0xf27f140   (G-copy)
  └─ FirmwareEventBuilder<…vlc>::FirmwareEventBuilder(   @0xf284d20
         TpuXPlaneBuilder*, d,d,d,d, m,m,m,m)
       ├─ doubles  ← row[+0x360 / +0x368 / +0x370 / +0x378]   (XMM args)
       ├─ ulongs   ← row[+0x380 / +0x388 / +0x390 / +0x398]   (GP args)
       └─ stamps "power" / "temperature" XStats per firmware event

constructor signature 由其 mangling 在字节层面确认：…FirmwareEventBuilder…C2EPNS_16TpuXPlaneBuilderEddddmmmm，即一个 TpuXPlaneBuilder* 后跟四个 double 和四个 unsigned long。reader 将四个 double 从 row[+0x360..+0x378] 加载到 xmm0..xmm3，将四个 ulong 从 row[+0x380..+0x398] 加载到 GP argument register，然后构造 builder。builder 将它们存入 obj+0x10..+0x48，并在 power/temperature 公式中消费它们。每个 coefficient 的精确角色（voltage²、current、thermal slope；ulong 作为 P-state id 或 thermal threshold）这里未解码；builder 内部见 v7x Perf-Counters § Firmware Event Model。

已消费偏移普查 ​

结合上面四个 reader 与 geometry/clock reader，表可划分为 有二进制内 reader 的字段 和 仅由 host roofline/power model 读取的字段。后者是此二进制盖印但不解释的 spec 常量。

由 libtpu 读取（CONFIRMED-CONSUMED） ​

Offset(s)字段Reader
+0x00core/multi flag（BYTE）plane setup, GetJobInfoFromResponse
+0x04 / +0x50 / +0x2f8GTC / TensorCore / SparseCore clock（kHz）GetJobInfoFromResponse（×1000 → Hz）
+0x0c / +0x10 / +0x14cores/chip、logical-device countToDeviceOrdinal（÷）、HandleConvolution、ExtractUtilizationCounters
+0x60 / +0x78 / +0x80peak FLOP/s bf16 / int8 / int4HandleConvolution（vminsd）、V2 stat 0x62
+0xd0peak HBM bandwidth（GB/s）V2 stat 0x63（×GiB）
+0x100 / +0x108 / +0x120 / +0x128per-mem-space bandwidth（GB/s）V2 stats 0x66/0x67/0x64/0x65（×GiB）
+0x2c4has_megacore（BYTE）V2 stat 0x6b、GetCostAdjustmentFunction
+0x2c8 / +0x348 / +0x350 / +0x358 / +0x438 / +0x440perf-counter-set baseGetPerformanceCounterNames<28/12/3>
+0x360..+0x378 (d) / +0x380..+0x398 (m)power/thermal bundleFirmwareEventBuilder
+0x2b8 / +0x2f8 / +0x340packed geom、SC clock、SC lane countProcessCounter、SC subscriber

仅在二进制外读取（INFERRED — 无 in-libtpu reader） ​

text
+0x28, +0x2d0   two 8-point DVFS frequency ladders (v7x-only)
+0x58, +0x68    per-LD bf16 / v7x int8-alt peak
+0xb8..+0xc8    effective HBM bandwidth class (value-scaled to bw; authoritative bw is +0xd0)
+0xd8..+0xf0    memory-latency class
+0xf8/+0x110/+0x118/+0x130   extra bandwidth/count members (not stamped this build)
+0x138..+0x150  secondary rate "B"
+0x178..+0x190  secondary rate "C"
+0x300/+0x308   core voltage / power
+0x438/+0x440 tail — verified perf-counter-set bases above, NOT a roofline double
```text

在 `.text` 中没有任何 `[base + ordinal*0x448 + disp]` load 使用这些 displacement（94-site reference census）。它们是 host xprof roofline/utilization/power model 的输入，被盖印到 XPlane 上（bandwidth/peak 子集）或作为 host tool 再选择的 spec 携带。

> **怪癖 —** 对“roofline 在哪里？”的回答是：libtpu 是 **roofline 输入的 producer，而不是 evaluator**。它把 `peak_teraflops_per_second` + 四个 `peak_*_bw_gigabytes_per_second` stat + `has_megacore` + perf-counter set + power/thermal bundle 盖印到 device XPlane 上；TFLOP/s-ceiling-vs-effective-bandwidth classifier 运行在读取该 plane 的 host xprof 工具中。若重新实现者寻找一个在 libtpu 内对 `+0xb8/+0xd8/+0x138` 执行 roofline 算术的函数，是找不到的；这些 double 未经读取就离开了二进制。

---

## 相关组件

| 组件 | 关系 |
|---|---|
| `GetJobInfoFromResponse` @ `0xf2c9ac0` | 读取 `+0x00/+0x04/+0x50/+0x2f8`，将三个 clock `×1000` kHz→Hz 提升到 device-info record |
| `ToDeviceOrdinal` @ `0xf69c580` | 读取 `+0x0c/+0x10/+0x14` geometry，以映射 `(host, core)` → device ordinal |
| `GetCostAdjustmentFunction` @ `0xf58efc0` | 读取 `+0x2c4` megacore flag + `+0x0c` cores/chip，以选择 cost-adjust function |
| `ExtractUtilizationCounters` @ `0xf0f7fa0` | 读取 `+0x0c` 作为 utilization denominator 的 `double` divisor |
| `FirmwareEventBuilder` ctor @ `0xf284d20` | 保存 `+0x360..+0x398` bundle；发射 `power`/`temperature` |

## 交叉引用

- [按 DeviceType 的 Profiler 结构](per-devicetype-struct.md) — 姊妹页面；本页读取的字节级 `0x448` 字段布局和 `DeviceType`↔codename 映射
- [v7x Perf-Counters](v7x-perf-counters.md) — 负责 `+0x2c8/+0x348/+0x350/+0x358/+0x438/+0x440` descriptor 字段、`DeviceType == 12` gate 和 `FirmwareEventBuilder` event model
- [Profiling 与 Telemetry](overview.md) — 该 stamping 馈入的 XPlane/XStat pipeline
- [Task Proto](task-proto.md) — 从 `+0x04/+0x50/+0x2f8` 填充的 device-info clock 字段（`gtc/tensor_core/sparse_core_freq_hz`）
- [TpuHloCostAnalysis](../cost/tpu-hlo-cost-analysis.md) — 共享 `HandleConvolution` peak-FLOP/s 读取的 cost model
- [Memory Bandwidth & Latency Model](../cost/memory-bandwidth-latency-model.md) — bandwidth 侧 cost 路径；host-roofline consumer frame
- [LocalDmaBandwidth](../cost/local-dma-bandwidth.md) — 与 `+0xb8/+0xd0` 平行的 compiler 侧按世代 bandwidth accessor
- [chip_parts.binarypb](../targets/chip-parts-binarypb.md) — *另一个* 按世代 spec 来源；此处用 `FlopsPerSecond` 交叉验证 `+0x60/+0x78/+0x80`