Internal Wiki

主题

Per-DeviceType Profiler 结构体 ​

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

摘要 ​

profiler 在一个烘入的数组中为每个硅片世代保存一个描述符 — xprof::kDeviceTypeInfo，共 17 条记录，步长 0x448（1096 字节），由 xprof::DeviceType 枚举序号直接索引。每条记录都是每个 DeviceType 的 profiler 结构体：它携带 trace pipeline 用来做除法的两个时钟（+0x04 处的 GTC 时钟和 +0x50 处的 TensorCore/compute 时钟）、+0x08 处的 GTC 时间戳位宽、+0x0c..+0x20 处的每芯片 core 几何、+0x28 和 +0x2d0 处的两条 8 点 DVFS 频率阶梯，以及大约四十个 IEEE-754 硬件规格 double。本页负责该结构体的字段布局和代号绑定 — 捕获到的 TPU 的 PCI 身份如何选择正确的 DeviceType 序号，该序号继而选择这个结构体实例，并进一步驱动 trace codec 的时钟域。兄弟页面 kDeviceTypeInfo 生产者和 Roofline 读取者 负责正交问题，即谁读取每个字段；本页负责每个字段是什么，以及每一行描述哪个世代。

最重要的结构事实是，该结构体没有指针。完整的 .rela.dyn 扫描在 [0x1c60480, 0x1c64d48) 内返回零个重定位 — 每个字段都是链接时冻结的内联 int32、int32[8] 或 double 字面量。这意味着每世代代号字符串和每世代 trace-codec 工厂没有嵌入这个结构体。它们是两个并行查找，由同一个捕获设备身份作为键，但存储在别处：代号字符串位于 0x21772f00 的 xprof::DeviceTypeString 指针数组中，codec 工厂位于通过 GetTraceCodec（0xf5a2900）到达的每家族 std::map 中。DeviceType 序号选择时钟（这个结构体）；原始 DeviceIdentifiers PCI 元组选择codec；二者在 trace 转换时汇合。若重实现者预期这个结构体持有代号或 codec 函数指针，就会寻找一个并不存在的字段。

绑定链是一个三跳 join，全部锚定到反汇编函数。捕获 trace 的 DeviceIdentifiers（一个 12 字节 PCI 元组）输入 DeviceTypeFromDeviceIdentifiers（0xf6993a0），该函数逐字段将元组与 kXxxChipIdentifiers 常量比较，并产生 {3,5,7,8,10,11,12,13} 中的一个 DeviceType 序号 — 八个真实硅片世代。该序号输入 GtcSpanConverter::GtcSpanConverter(DeviceType)（0xf2cb6e0），后者加载 kDeviceTypeInfo[ordinal] + 0x04（GTC kHz 除数），同时输入 DeviceTypeString(DeviceType)（0xf69c7c0），后者产出设备平面名称（"TPU v2" … "TPU v7x"）。本页走查这条链，解码字段偏移图，并将每个序号固定到其 codec 家族代号（jxc/pxc/vxc/gxc 以及每世代的 pfc/plc/vfc/vlc/glc/gfc 子家族）。

对重实现而言，契约是：

  • 0x448 字节字段偏移图：两个时钟（+0x04 GTC kHz，+0x50 compute kHz）、+0x08 GTC 时间戳宽度、+0x0c..+0x20 几何、+0x28/+0x2d0 DVFS 阶梯，以及每世代规格 double 区域 — 将确定的头部字段与推断的大块字段分开。
  • DeviceType 序号 → 结构体实例寻址：base + ordinal*0x448、ordinal >= 0x11（17 行）边界，以及 0x1c60480 和 0x1c84d90 处两个字节完全相同的 ICF 副本。
  • 代号绑定：DeviceIdentifiers PCI 元组 → DeviceType 序号（DeviceTypeFromDeviceIdentifiers）→ {clock, name}，并行地 → codec（GetTraceCodec）；8 世代总表。
  • “无指针字段”不变量：结构体内零重定位 ⇒ 代号和 codec 单独键控，而非嵌入。
结构体符号xprof::kDeviceTypeInfo (_ZN5xprofL15kDeviceTypeInfoE) @ 0x1c60480 (.lrodata)
布局17 条记录 × 0x448 = 18632 字节；结束于 0x1c64d48（下一个符号 xla::gpu::kDeviceHloOpProfiles @ 0x1c64d50）
第二副本字节完全相同的 ICF fold @ 0x1c84d90（由 NormalizeGtcTimestamps 读取）
索引xprof::DeviceType 序号（0..16）；8 个真实世代 = {3,5,7,8,10,11,12,13}
时钟消费者GtcSpanConverter::GtcSpanConverter(DeviceType) @ 0xf2cb6e0（只读取 +0x04）
代号绑定DeviceTypeFromDeviceIdentifiers @ 0xf6993a0；DeviceTypeString @ 0xf69c7c0
Codec 绑定（并行）GetTraceCodec @ 0xf5a2900（由 DeviceIdentifiers 键控，不由此结构体键控）
指针字段无 — [0x1c60480, 0x1c64d48) 内零重定位

结构体 — 布局和寻址 ​

目的 ​

kDeviceTypeInfo 是 profiler 的每世代描述符表。每个 DeviceType 序号对应一条 0x448 字节记录，保存 trace pipeline 需要了解、但 trace 流自身不包含的关于某个硅片世代的一切：必须用来除原始计数器 tick 的时钟频率、必须用于掩码的时间戳计数器宽度、core 几何，以及一块每世代硬件规格 double。到达一条记录只需一次乘法和边界检查：序号乘以 0x448 步长，并由 17 行检查保护。

入口点 ​

text
DeviceIdentifiers (PCI tuple, from the captured JfTrace)
  └─ DeviceTypeFromDeviceIdentifiers (0xf6993a0)   ── PCI tuple → DeviceType ordinal
       └─ GtcSpanConverter::GtcSpanConverter(ordinal) (0xf2cb6e0)
            └─ row = kDeviceTypeInfo + ordinal*0x448 ; clk = row[+0x04]
```text

### 寻址

每个读取者都以同样方式到达一条记录：相对 GOT 物化表基址（这样代码可针对链接器分配的任一 ICF 副本工作），以 17 为序号做边界检查，再乘以 `0x448` 步长。规范加载点是 `GtcSpanConverter` 构造函数，它恰好读取一个字段。

```c
function GtcSpanConverter_ctor(GtcSpanConverter* this, uint ordinal):   // 0xf2cb6e0
    if ordinal >= 0x11:                  // 17-row bound (ud1 trap)
        trap()
    // 274 int32 elements == 0x448 bytes; the trailing offset lands on +0x04
    this[0] = *(int32*)(kDeviceTypeInfo + 274*ordinal + 1)   // == row[+0x04], GTC kHz
    this[1] = 0                                              // GtcSpan vector cleared
    this[2] = 0
    this[3] = 0

注意 — IDA 反编译器把行步长呈现为 int* 上的 274 * ordinal（因此 274 × 4 = 0x448 字节），而常量折叠 + 274*ordinal − 8525971 解析为 kDeviceTypeInfo + ordinal*0x448 + 4。不要把 274 读作字段数量；它是用 int32 元素表达的 0x448 字节步长。相同的 >= 0x11 guard 出现在每个加载点。

两个 ICF 副本 ​

nm 在多个 VMA 处列出 _ZN5xprofL15kDeviceTypeInfoE（以及一串带 _0、_1、… 后缀的 fold）。其中两个是本页讨论的 live 表：0x1c60480 处的主表（由 GtcSpanConverter 读取）和 0x1c84d90 处字节完全相同的 fold（由 NormalizeGtcTimestamps<...> 家族读取）。两个副本的时钟列已在全部 17 行中逐 dword 比较，结果完全相同。mangling 中的 L 标记内部链接 — 一个翻译单元局部的 static const 数组 — 并且由于它被许多模板/pass 实例化引用，链接器的 identical-code-folding pass 复制了数据，而不是把它合并为一个符号。这些副本是 folding 产物，不是独立表。

易错点 — NormalizeGtcTimestamps<...> 实例化读取的是第二副本（0x1c84d90），不是主表。NormalizeGtcTimestamps<pxc::profiler::TraceEntry>（0xf59b080）加载 *(int*)(GLOBAL_OFFSET_TABLE_ + 274*ordinal − 136378107)，它解析为 0x1c84d90 + ordinal*0x448 + 4 — 同一个 +0x04 GTC-clock 字段，但来自 folded 副本。重实现中只需要一张表；双副本拆分是链接期细节，两个副本被读取的目的相同（+0x04 时钟）。

字段偏移图 ​

目的 ​

这是对 0x448 记录的字节级解码。只有两个字段有直接追踪到的、在二进制中读取这个具名结构体且处于时钟语境的消费者 — +0x04（GTC 时钟，由 GtcSpanConverter 和 NormalizeGtcTimestamps 读取）以及通过值交叉检查确认的 +0x08（GTC 宽度，与 codec 的 GetBits64 宽度逐字节匹配）。两个时钟以下的所有内容都是字节精确且按世代单调的，但其成员名称是根据值缩放推断的，因为消费这些字段的 libtpu 内读取者（在兄弟页面中编目）通过原始位移访问，而非具名访问器。置信度也相应标注。

头部标量和时钟 ​

字段偏移类型含义
core_multi_flag+0x00int32 (BYTE-used)在 2-core / mega / SparseCore 世代（DT 3,5,7,8,10,11,12,13）上为 1，在单 core 占位槽上为 0
gtc_freq_khz+0x04int32 (kHz)GTC (Global Time Counter) 频率。皮秒时间基准除以 (khz << 4)。GtcSpanConverter 读取的唯一字段。
gtc_ts_width_bits+0x08int32 (bits)GTC 时间戳计数器宽度 {48, 45, 64}。与每家族 codec GetBits64 宽度逐字节匹配；wrap 周期 = 2^width / (khz·1000) s
cores_per_chip+0x0cint32在 megacore 世代（DT 3,5,7,10,12）上为 2，其他为 1 — TensorCore 数量
geom_a+0x10int32pre-SparseCore 世代上为 {1,2,4}，SC 世代上为 0 — TensorCore 侧 lane/MXU group 数
geom_b+0x14int32只在 SC 世代非零（DT10=4, DT12=4, DT13=2） — SparseCore 侧计数
geom_c+0x18int32多数为 1；DT3/DT5 上为 2 — 每芯片乘数
geom_d+0x1cint32递增的 {1,2,4,6,8}（DT7=4, DT9=2, DT10=6, DT12=8） — 每世代 tile/engine 数
sc_present_flag+0x20int32只在 45-bit SC 世代（DT10..13）上为 1，其他为 0
dvfs_ladder_1+0x28int32[8] (kHz)8 点 DVFS 工作点阶梯；在 DT12 上填充（{1600000..2200000}）
compute_clk_khz+0x50int32 (kHz)TensorCore/compute 时钟，不同于 +0x04 处的 GTC 时钟；按世代递增

特性 — GTC 时钟（+0x04）和 compute 时钟（+0x50）在多数世代上是不同频率，绝不能混淆。DT12（v7x）运行 833000 kHz GTC 计数器，但 compute 时钟为 1900000 kHz；DT5（v3）的 GTC 时钟为 700000 kHz，compute 时钟为 940000 kHz。trace 时间戳处于 GTC 域（+0x04）；compute 时钟（+0x50）是 cost model 使用的 cycle 时钟。用 compute 时钟给时间戳定标的重实现会使每个事件产生偏斜。

规格 Double 和 DVFS 区域 ​

+0x50 以下，记录主要由 IEEE-754 double 和第二条 DVFS 阶梯组成。这些是每世代硬件规格常量 — peak compute、memory bandwidth、latency 和 DVFS/voltage class — 按行冻结。它们的精确成员身份无法从此二进制恢复（没有 ToString，profiler-clock 语境中没有逐字段具名访问器），因此这些区域按其形状和每世代缩放描述，而不是逐行转录。

区域偏移范围类型特征（每世代缩放）
Compute-class doubles+0x58 .. +0x80double单调递增的 peak-compute 指标；+0x60 跟踪类似 TFLOP/s 的数值（24.3 DT3 → 1029 DT12），+0x78/+0x80 约为 +0x60 的 2×/4×（精度层级）
Bandwidth-class doubles+0xb8 .. +0xd0double类 GB/s 带宽；+0xb8 280(DT3) → 3433(DT12)；权威 HBM 数值是 +0xd0
Latency-class doubles+0xd8 .. +0xf0doublelatency/cycle 类；+0xe8==+0xe0 且 +0xf0==+0xd8（每世代重复项）
Count-class doubles+0xf8 .. +0x130double大计数（peak-ops / systolic-cell scale）；每精度变体
Secondary-rate groups+0x138 .. +0x194double仅 DT7+（56.6、453 家族位于 +0x13c/+0x17c）
packed_geom+0x2b8int32打包的 {a,b,a,b} 4 字节几何描述符（例如 DT3 = 01 08 01 08）
has_megacore-class+0x2c4int32 (BYTE)在 DT7/DT10 上为 1 — megacore 风格标志
Perf-counter-set bases+0x2c8 / +0x348 / +0x350 / +0x358 / +0x438 / +0x440int32v7x perf-counter set 的打包 enum base（仅在 DT12 上非零） — 归 v7x Perf-Counters 所有
dvfs_ladder_2+0x2d0int32[8] (kHz)第二条 8 点 DVFS 阶梯；DT12 = {1400000..1900000}
dvfs_nominal_khz+0x2f8int32 (kHz)标称 DVFS / SparseCore 工作点（1750000 DT12，pre-SC 世代为 0）
Voltage/power doubles+0x300 / +0x308doublevoltage/power 类（3.6/5.85 DT12）
sc_lane_count+0x340int32SC 世代（DT10/12/13）上为 16，其他为 0
Firmware calib bundle+0x360 .. +0x398double×4 + ulong×4输入 FirmwareEventBuilder 的每世代 power/thermal 系数

注意 — +0x438/+0x440 尾部保存两个 int32 perf-counter-set enum base — 用于 v7x ICR（+0x438）和 CMNUR/HBM（+0x440）counter set，与 v7x Perf-Counters 页面从 0x1c637e0 的 DT12 行恢复出的六个描述符字段相同。它们只在 DT12 上非零，不包含指针（高 dword 为零），也不是 roofline double。resolver 调用点确认了加法式 base + ordinal*8 用法。

皮秒时间基准 ​

+0x04 时钟由 GtcSpanConverter::TimespanFromGtcSpan（0xf2cb7e0）作为 GTC→host-ns 转换的除数消费。该算术是精确的，并固定了该字段的角色和单位。

c
function TimespanFromGtcSpan(this, gtc_ticks):                 // 0xf2cb7e0
    khz   = this[0]                          // == kDeviceTypeInfo[ord]+0x04
    scale = 16 * khz                         // the GtcSpan ×16 fixed-point (low 4 bits fractional)
    // value_ns = round( gtc_x16 * 1e9 / (khz << 4) )
    numer = (scale >> 1) + 1000000000 * (gtc_ticks - anchor)   // 0x3B9ACA00 == 1e9
    ns    = udiv128(numer, scale)
    return ns + 1000 * base_ns               // bracketed by the GtcSpan rb-tree node
```text

> **注意 —** 单位证明是字节算术。由于 GtcSpan 携带 ×16 定点值，`value_ps = gtc_x16 · 1e9 / (khz << 4)`。对于一个整数 tick（`gtc_x16 = 16`）：`700000 → 1429 ps`（`1/(700·1e6)·1e12 = 1428.57`）、`800000 → 1250`、`833000 → 1200`、`1333000 → 750` — 每个都与 `1/(khz·1000)·1e12` 在 1 ps 内匹配。GTC wrap 周期由 `2^width / (khz·1000)` 得出：48-bit/700 MHz 计数器约每 111.7 h wrap 一次；45-bit/800 MHz 计数器约每 12.2 h wrap 一次；`DT9` 的 64-bit/1.333 GHz 槽实际不会 wrap。因此 `+0x04` 时钟是*权威* GTC 除数，按世代烘入 — 不是运行时 `Task.gtc_freq_hz` 字段。

---

## DeviceType → 代号绑定

### 目的

此结构体由 `DeviceType` 序号索引，但序号本身来自捕获硬件的 PCI 身份。该绑定是一组函数：将 `DeviceIdentifiers` 元组映射到序号并可反向映射，同时将序号映射到公开名称。它们共同把 `kDeviceTypeInfo` 的每一行固定到具体硅片世代及其 codec 家族代号。

### 入口点

```text
DeviceTypeFromDeviceIdentifiers (0xf6993a0)   ── PCI tuple → DeviceType ordinal (StatusOr)
  ├─ field-compare vs jxc::kJellyfishIdentifiers, kDragonfishIdentifiers
  ├─ field-compare vs pxc::plc::kPuffyliteChipIdentifiers
  ├─ field-compare vs pxc::pfc::kPufferfishChipB0{Mfg,Water,Air}Identifiers
  ├─ field-compare vs vxc::vlc::kViperliteChip{A0,A1}{VF,PF}Identifiers
  ├─ field-compare vs vxc::vfc::kViperfishChip{VF,PF}Identifiers
  ├─ IsGlc (0xf6992a0)  → ordinal 13   ── gxc::glc (Ghostlite)
  ├─ IsGfc (0xf699320)  → ordinal 12   ── gxc::gfc (Ghostfish)
  └─ default → MakeErrorImpl<3>("Unsupported device identifiers")
DeviceTypeString (0xf69c7c0)                  ── ordinal → public name array @0x21772f00
DeviceIdentifiersFromDeviceType (0xf6996e0)   ── ordinal → representative PCI tuple (inverse)

算法 ​

DeviceTypeFromDeviceIdentifiers 逐字段将捕获的 12 字节元组与 kXxxChipIdentifiers 常量比较，并存储匹配的序号。Ghostlite/Ghostfish 家族通过 IsGlc/IsGfc 谓词分派，而不是内联常量（每个谓词覆盖多个 App/Mgt SKU）。无匹配时产生错误码 3。

c
function DeviceTypeFromDeviceIdentifiers(out, const DeviceIdentifiers* id):  // 0xf6993a0
    if id == jxc::kJellyfishIdentifiers:          out.ordinal = 3;  return ok   // TPU v2
    if id == jxc::kDragonfishIdentifiers:         out.ordinal = 5;  return ok   // TPU v3
    if id == pxc::plc::kPuffyliteChipIdentifiers: out.ordinal = 8;  return ok   // TPU v4 Lite
    if id in pxc::pfc::kPufferfishChipB0{Mfg,Water,Air}:
                                                  out.ordinal = 7;  return ok   // TPU v4
    if id in vxc::vlc::kViperliteChip{A0,A1}{VF,PF}:
                                                  out.ordinal = 11; return ok   // TPU v5 Lite
    if id in vxc::vfc::kViperfishChip{VF,PF}:      out.ordinal = 10; return ok   // TPU v5
    if IsGlc(id):                                 out.ordinal = 13; return ok   // TPU v6 Lite
    if IsGfc(id):                                 out.ordinal = 12; return ok   // TPU v7x
    return MakeErrorImpl<3>("Unsupported device identifiers")  // device_identifiers_utils.cc:152
```text

`DeviceTypeString` 是 ordinal→name 映射：`ord-1` 索引 `0x21772f00` 处的指针数组（有效 `0..0xC`，即序号 `1..13`）；任何其他序号返回常量 `"Cloud TPU"`。`DeviceTypeToHardwareType`（`0xf69c7a0`）是并行的 ordinal→hardware-class 映射（`0xAB8A2F4` 处的 int32 数组）：GPU=2，TPU=3，占位符为 0/1。

```c
function DeviceTypeString(int ordinal):           // 0xf69c7c0
    if (uint)(ordinal - 1) > 0xC:  return "Cloud TPU"
    return off_21772F00[ordinal - 1]              // "GPU","TPU v2",...,"TPU v6 Lite"

总表 ​

八个真实硅片世代正是 DeviceTypeFromDeviceIdentifiers 能返回的序号（全部 hardware-type 3 = TPU）。其余序号是 GPU/占位/保留槽，携带默认时钟但没有代号也没有 PCI 元组。下方的公开名称、codec 家族、时钟和时间戳宽度均已通过 DeviceTypeString、DeviceTypeFromDeviceIdentifiers 中的 codec namespace，以及直接读取该结构体 +0x04/+0x08 列逐字节确认。

DT公开名称代号Codec 家族GTC kHz (+0x04)ts-width (+0x08)compute kHz (+0x50)hwtype
1GPU(host GPU plane)—700000487000002 (GPU)
2Cloud TPU(generic placeholder)—700000487000000
3TPU v2Jellyfishjxc700000487000003 (TPU)
4Cloud TPU(placeholder; bind → err)—700000487000001
5TPU v3Dragonfishjxc700000489400003
6Cloud TPU(generic placeholder)—700000487000000
7TPU v4Pufferfishpxc::pfc7000004810500003
8TPU v4 LitePuffylitepxc::plc7000004810500003
9Cloud TPU(reserved 64-bit slot)—13330006413330000
10TPU v5Viperfishvxc::vfc8000004517500003
11TPU v5 LiteViperlitevxc::vlc8000004515000003
12TPU v7xGhostfishgxc::gfc8330004519000003
13TPU v6 LiteGhostlitegxc::glc8000004517500003
14–16Cloud TPU(legacy placeholders)—70000048700000—

特性 — DeviceType 12（"TPU v7x"，Ghostfish/gfc）和 DeviceType 13（"TPU v6 Lite"，Ghostlite/glc）是同一个 gxc 芯片家族的兄弟成员，但它们是不同的 DeviceType，具有不同的 GTC 时钟（833 vs 800 MHz）和不同的 compute 时钟（1.9 vs 1.75 GHz）。二者都通过谓词函数（IsGfc/IsGlc）分派，而不是单个常量，因为各自跨越多个 App/Mgt PCI SKU。把整个 gxc 家族当成一个 DeviceType 的重实现会为其中一半选择错误的时钟除数。

易错点 — DeviceType 9 是保留槽：1333000 kHz GTC 时钟、64-bit（实际不会 wrap）时间戳、名称为 "Cloud TPU"，且没有 PCI 元组 — DeviceIdentifiersFromDeviceType(9) 返回错误。它不是八个具名世代之一；其时钟/宽度是唯一的非占位字段。不要把它绑定到 codec；此构建中没有任何代码这样做。

反向映射和 PCI 元组 ​

DeviceIdentifiersFromDeviceType（0xf6996e0）是反向映射：对序号做 switch（case 3,5,7,8,10,11,12,13），返回一个代表性的 kXxxChipIdentifiers 元组；序号 4/6/9（以及 switch 外的任何值）返回 MakeErrorImpl<3>("Unsupported device type")。这些元组是 12 字节 PCI 描述符 — +0 vendor、+2 device、+4 subsys_vendor、+6 subsys_device、+8 class、+9 subclass、+0xa prog_if、+0xb revision — 全部具有 vendor_id == 0x1AE0（Google）和 subsys_vendor == 0x1AE0。在 0xbdf3c0c..0xbdf3cdc 中有十七个此类元组（最后一个 gfc 从 0xbdf3cd0 开始并在 0xbdf3cdc 结束）；它们几乎连续，只被两个 jxc 元组之后的 4 字节对齐空隙（0xbdf3c18→0xbdf3c28）打断。只有 kPuffyliteChipIdentifiers @ 0xbdf3c4c 带有 ELF 符号（nm 恰好报告一个 *ChipIdentifiers 符号）；其他十六个元组 — 包括 0xbdf3cc4/0xbdf3cd0 处两个 Ghostfish（gfc）元组 — 在符号表中未命名，每个都由其 Is<Family> 访问器的内联比较识别（例如 IsGfc）。DeviceTypeFromDeviceIdentifiers 中的字段比较会掩码比较，使 WORD2 == 6880（0x1AE0）subsystem-vendor 检查和 revision-byte 检查共同 gate 每个匹配。

c
function DeviceIdentifiersFromDeviceType(out, int ordinal):   // 0xf6996e0
    switch ordinal:
        case 3:  tuple = jxc::kJellyfishIdentifiers              // TPU v2
        case 5:  tuple = jxc::kDragonfishIdentifiers             // TPU v3
        case 7:  tuple = pxc::pfc::kPufferfishChipB0WaterIdentifiers
        case 8:  tuple = pxc::plc::kPuffyliteChipIdentifiers
        case 10: tuple = vxc::vfc::kViperfishChipVFIdentifiers
        case 11: tuple = vxc::vlc::kViperliteChipA0VFIdentifiers
        case 12: tuple = <gfc tuple @0xbdf3cd0>                  // device 0x76 / subsys 0xf2
        case 13: tuple = gxc::glc::kGhostliteChipAppVFIdentifiers
        default: return MakeErrorImpl<3>("Unsupported device type")  // ...utils.cc:191
    out = { ok, tuple.lo64, tuple.dword2 }
```text

---

## 为什么结构体不携带代号或 Codec 指针

对 `r_offset` 位于 `[0x1c60480, 0x1c64d48)` 内的任何重定位进行 `.rela.dyn` 扫描，返回**零**。`0x448` 记录是纯标量/浮点数据；其中没有任何已重定位指针。由此得到两个后果，二者在上文已经使用：

- **代号字符串**不在结构体中。它位于 `0x21772f00` 的 `DeviceTypeString` 指针数组中（在 `.data.rel.ro`，其中指针*会*被重定位），由同一个 `DeviceType` 序号索引。该序号把时钟（此结构体）join 到名称（那个数组）。
- **trace-codec 工厂**不在结构体中。它注册在一个由 12 字节 `DeviceIdentifiers` 值键控的每家族 `std::map` 中，并通过 `GetTraceCodec`（`0xf5a2900`）到达。`DeviceType` 序号选择*时钟*；原始 PCI 元组选择*codec*。它们是两个由同一捕获设备身份键控的并行查找，在 trace 转换时 join。

因此，从捕获 trace 到解码、时间转换并命名的事件的完整主桥梁是：

```text
captured JfTrace.device_identifiers (12-byte PCI tuple)
  ├─ GetTraceCodec(tuple)                    → per-gen TraceCodecInterface (the packet decoder)
  └─ DeviceTypeFromDeviceIdentifiers(tuple)  → DeviceType ordinal
       ├─ GtcSpanConverter(ordinal)          → kDeviceTypeInfo[ordinal]+0x04 (GTC ps timebase)
       └─ DeviceTypeString(ordinal)          → device-plane name ("TPU vN")

注意 — 并行查找设计正是该结构体无指针的原因。如果 codec 工厂或代号被嵌入，结构体就会携带重定位，并且无法成为 .lrodata 中的纯 static const 聚合。保持时钟（由序号键控）和 codec（由元组键控）分离，使表保持无重定位，并允许同一个 DeviceType 序号作为 join key 进入三个独立数组。重实现可以把代号存进结构体，但二进制有意不这样做 — 而“无重定位”不变量就是证明。

未决问题 ​

这些字段和行为是字节精确的，但其语义或消费者无法仅从此二进制固定。

  • 大块字段成员名称。 约 40 个规格 double（+0x58..+0x194、+0x300/+0x308、+0x360..）和整数几何（+0x0c..+0x20、+0x2b8、+0x340、+0x388..）按世代单调（compute / bandwidth / DVFS / count 类），但在 profiler-clock 路径中没有具名访问器；其成员身份由值缩放推断。兄弟读取者页面 追踪 stat-stamp 和 cost-model 路径消费哪些位移。
  • codec 是否读取 +0x08。 GTC 宽度 {48/45/64} 与每家族 GetBits64 codec 常量逐字节完全匹配，但 codec 的宽度看起来烘入在家族模板化的 DecodeEntry 中，而不是运行时从 kDeviceTypeInfo[ord]+0x08 加载 — 因此 +0x08 可能是权威但仅供信息参考的副本，而不是 live source。
  • DeviceType 9 和 14–16。 保留的 64-bit/1.333 GHz 槽（DT9）以及三个 700 MHz/48-bit 占位符（DT14–16）携带时钟和 pre-SC 几何，但没有代号也没有 PCI 元组；它们预期对应哪种硅片（如果有）无法在此恢复。

相关组件 ​

组件地址关系
GtcSpanConverter::GtcSpanConverter(DeviceType)0xf2cb6e0规范读取者；加载 kDeviceTypeInfo[ord]+0x04 作为 GTC ps 除数
GtcSpanConverter::TimespanFromGtcSpan0xf2cb7e0使用 +0x04 时钟应用 ×16 / 1e9 ps 转换
NormalizeGtcTimestamps<...>0xf59b080 (pxc)为同一时钟读取第二个 ICF 副本 0x1c84d90 + ord*0x448 + 4
DeviceTypeFromDeviceIdentifiers0xf6993a0PCI 元组 → DeviceType 序号（绑定到此结构体索引）
DeviceIdentifiersFromDeviceType0xf6996e0反向：序号 → 代表性 PCI 元组
DeviceTypeString0xf69c7c0序号 → 公开名称（0x21772f00 ptr array）— 代号的独立位置
DeviceTypeToHardwareType0xf69c7a0序号 → 硬件类别（0xAB8A2F4 int32 array）
GetTraceCodec0xf5a2900并行的元组键控 codec 查找 — 选择 decoder，而非时钟
IsGlc / IsGfc0xf6992a0 / 0xf699320gxc::glc（DT13）和 gxc::gfc（DT12）家族的谓词分派

交叉引用 ​