逐代精度表

本页所有值、blob offset 和常量，均从五个嵌入式 accuracy_table_*.binarypb blob 以及来自 libtpu-0.0.40-cp314 wheel 的 libtpu.so 中的 select_accuracy.cc consumer function 逐字节精确解码得到（BuildID md5 89edbbe81c5b328a958fe628a9f2207d）。其他版本会有所不同。

摘要

每个 TPU generation 都携带一个嵌入式 accuracy_table_<codename>.binarypb proto blob，用来记录 硬件 transcendental 近似的数值精度：sin、cos、tan、exp、log、rsqrt、sqrt、divide、erf、logistic、tanh 以及它们的 fast/accurate 变体。这些不是原始 minimax polynomial coefficient；它们是 测得的 最坏情况 relative-error 和 ULP bound，加上一条分段 input→error 曲线，编译器会查询它们来决定 fast hardware approximation（XLU/EUP 路径）是否满足用户请求的 ResultAccuracy，还是必须发射更慢的 accurate software polynomial。

共有五个 blob：jellyfish (TPU v2)、dragonfish (TPU v3)、pufferfish (TPU v4)、viperfish (TPU v5)、ghostlite (TPU v6 lite)；外部名称按 TpuVersionToExternalName（见 codename matrix）。最新 generation 6acc60406（外部 TPU7x）没有 accuracy table：它的 embed:// lookup 会失败，impl-type selector 会落入由 warning 驱动的 fallback。（Trillium/Ironwood 是 marketing name，在二进制中出现次数为 零；6acc60406 是这一代唯一的内部名称。）前三代携带逐字节相同的内容（除了 version stamp 外，都是同一个 181,284-byte blob）；viperfish 和 ghostlite 都把 fast-path bound 收紧了一到四个数量级，并且与逐代 transcendental cycle-cost 的改善同步。

解码达到可重实现级别：全部五个 blob 都从 .lrodata 逐字节精确 carve 出来，并用 md5 与 FileWrapper fingerprint 验证；proto schema 从 message typeinfo 加 serializer wire tag 恢复；四函数 consumer chain（GetTargetAccuracy → GetAccuracyEntry → CompareAccuracy → GetImplType）已反汇编。

Blob	5 × accuracy_table_<codename>.binarypb in .lrodata, md5-verified vs FileWrapper fp
格式	raw uncompressed serialized protobuf (xla.service.jellyfish.accuracy.AccuracyTable)
每表 entry 数	19 (15 distinct op names; 6 ops carry a fast+accurate variant pair)
有表的 generation	jellyfish, dragonfish, pufferfish, viperfish, ghostlite — not 6acc60406
Version stamp	proto field 3 = tpu::TpuVersion + 1 (1..5)
Consumer	xla::jellyfish::{GetTargetAccuracy, GetAccuracyEntry, CompareAccuracy, GetImplType}
源文件	platforms/xla/service/jellyfish/select_accuracy.cc
置信度	CONFIRMED unless a cell is annotated otherwise

---

逐代 Blob 映射

每个 blob 都由其 embed:// URI 中的 codename 作为键，并带有 proto field 3 = TpuVersion + 1 的 stamp。两套版本编号同时存在，且不能混淆：运行时 tpu::TpuVersion（0-based，从 Target+0x398 读取）和 proto 内部 field-3 stamp（1-based）。

codename	tpu::TpuVersion	proto f3	blob file	data_va / off	size (B)	md5 (== FileWrapper fp) · Conf
jellyfish	0	1	accuracy_table_jellyfish.binarypb	0x040c8990	181,284	e3a2a768…d4c42cf4
dragonfish	1	2	accuracy_table_dragonfish.binarypb	0x040703d0	181,284	4d2efd10…748bc019
pufferfish	2	3	accuracy_table_pufferfish.binarypb	0x040f4dc0	181,284	e8dbf7b1…a185d4c4
viperfish	3	4	accuracy_table_viperfish.binarypb	0x041211f0	181,139	354963c9…3746b900
ghostlite	4	5	accuracy_table_ghostlite.binarypb	0x0409c800	180,617	8acb76b8…fb433612
6acc60406	5	—	— (none)	—	—	no table — fallback

embed:// URI 由路径片段 "embed://accuracy_table_" + codename + "_memfile/" + "accuracy_table_" + codename + ".binarypb" 组装成 embed://accuracy_table_<codename>_memfile/accuracy_table_<codename>.binarypb；codename 来自 tpu::TpuVersionToString(version)（string table @ 0x22011bf0，0→jellyfish … 5→6acc60406，大于 5 则 FATAL）。

GOTCHA — jellyfish、dragonfish 和 pufferfish 除 field-3 stamp 外逐字节相同（同为 181,284 B，同一组 per-op value，同一条 1,187-row exp curve），它们的 XLU/EUP 单元是同一设计。重实现者可以为 v2/v3/v4 发布一张表，只需要为 viperfish 和 ghostlite 提供不同数据。

---

Proto Schema

message AccuracyTable {                  // the blob
  repeated AccuracyEntry entries = 1;     // tag 0x0a; 19 per table
}
message AccuracyEntry {
  string name           = 1;  // 0x0a  "cosine","exponential",…
  PrimitiveType type    = 2;  // 0x10  always 11 == F32
  int32  version        = 3;  // 0x18  TpuVersion + 1
  int32  variant        = 4;  // 0x20  0 = fast/NoEup, 1 = EUP/accurate
  double max_input      = 5;  // 0x29  domain upper bound
  double worst_rel_err  = 6;  // 0x31  worst-case relative error over domain
  double worst_ulp      = 7;  // 0x39  worst-case ULP over domain
  repeated ErrorInfo curve = 8; // 0x42 piecewise input→error breakpoints
}
message AccuracyEntry.ErrorInfo {
  double input   = 1;  // 0x09  input-magnitude breakpoint
  double rel_err = 2;  // 0x11  relative error guaranteed at/below
  double ulp     = 3;  // 0x19  ULP bound at/below
}
```text

每个 entry 中的 `type`（field 2）都统一为 **11 = `F32`**：无论 bf16 I/O 如何，XLU 内部都以 F32 计算，因此同一个 entry 也管辖 bf16-cast 变体。`curve`（field 8）是单调的：随着 `|input|` 从 0/subnormal 增大，近似的 relative error 会缩小；`worst_rel_err`/`worst_ulp` 是最小输入 breakpoint 处的值。

19 个 entry（15 个 op name）是：三角函数 `cosine`、`sine`、`tan`；exp 家族 `exponential`（×2）、`exponential-minus-one`（×2）；log 家族 `log`（×2）、`log-plus-one`（×2）；根/除法 `rsqrt`、`sqrt`、`divide`；特殊函数/activation `erf`、`logistic`（×2）、`tanh`（×2）。六个双变体 op 同时携带 variant-0（fast / “NoEup” XLU）和 variant-1（高精度 “EUP”）entry。它们对应携带 `result_accuracy` attribute 的 StableHLO op（`vhlo.cosine_v2`、`vhlo.exponential_v2`，等等）。没有 `cbrt` entry（二进制字符串：`"Accuracy analysis for cbrt is not implemented yet."`）。

---

## 跨代精度常量

顶层最坏情况 **relative error**（field 6），按 `(op, variant)` 列出；越低表示越精确。`j/d/p` = jellyfish/dragonfish/pufferfish（三者相同）。**置信度：** 下列每个值均为 CONFIRMED，都是从 carve 出的 blob（[逐代 Blob 映射](#逐代-blob-映射) 中的 `data_va`）的 field-6/field-7 double 逐字节精确读取；每个单元格的来源是 proto，不是推断。

| op | var | jf/df/pf | viperfish | ghostlite |
|---|:---:|---:|---:|---:|
| cosine | 0 | 2.38e-07 | 2.38e-07 | 2.38e-07 |
| sine | 0 | 2.38e-07 | 2.38e-07 | 2.38e-07 |
| tan | 0 | 4.77e-07 | 4.77e-07 | 4.77e-07 |
| divide | 0 | 1.0 | 1.0 | 1.0 |
| rsqrt | 0 | 2.38e-07 | 2.38e-07 | 1.19e-07 |
| sqrt | 0 | 2.38e-07 | 2.38e-07 | 2.38e-07 |
| erf | 0 | 4.77e-07 | 9.54e-07 | 2.38e-07 |
| exponential | 0 / 1 | 1.0 / 1.0 | 1.0 / 1.0 | 1.0 / 1.0 |
| exp-minus-one | 0 / 1 | 2.44e-04 / 1.0 | 2.44e-04 / 1.0 | 1.0 / 1.0 |
| log | 0 / 1 | 4.88e-04 / 4.77e-07 | 3.82e-06 / 4.77e-07 | 2.38e-07 / 4.77e-07 |
| log-plus-one | 0 / 1 | 4.88e-04 / 1.19e-07 | 2.44e-04 / 1.19e-07 | 2.44e-04 / 1.19e-07 |
| logistic | 0 / 1 | 4.0 / 1.0 | 2.0 / 1.0 | 1.0 / 1.0 |
| tanh | 0 / 1 | 1.22e-04 / 4.77e-07 | 7.63e-06 / 4.77e-07 | 1.19e-07 / 4.77e-07 |

顶层最坏情况 **ULP**（field 7），布局相同：

| op | var | jf/df/pf | viperfish | ghostlite |
|---|:---:|---:|---:|---:|
| cosine | 0 | 4 | 4 | 4 |
| sine | 0 | 4 | 4 | 4 |
| tan | 0 | 8 | 8 | 8 |
| divide | 0 | 256 | 64 | 1 |
| rsqrt | 0 | 2 | 2 | 1 |
| sqrt | 0 | 4 | 4 | 2 |
| erf | 0 | 8 | 8 | 1 |
| exponential | 0 / 1 | inf / 1 | inf / 1 | inf / 1 |
| exp-minus-one | 0 / 1 | 2048 / 1.68e7 | 2048 / 1.68e7 | 1.68e7 / 1.68e7 |
| log | 0 / 1 | 4096 / 4 | 64 / 4 | 2 / 4 |
| log-plus-one | 0 / 1 | 4096 / 1 | 4096 / 1 | 4096 / 1 |
| logistic | 0 / 1 | 6.71e7 / 256 | 3.36e7 / 128 | 1.68e7 / inf |
| tanh | 0 / 1 | 2048 / 8 | 128 / 8 | 1 / 8 |

> **NOTE —** `exponential` variant-0 worst-ULP 在每个 generation 上都保持 `inf`，因为最小 breakpoint 位于 subnormal input，此处 relative error 为 1.0（100%）。fast `exp` 只在 subnormal range 以上才可用；`worst_*` 数值有意采用最坏情况（最小输入）值，而不是典型精度。

---

## 分段误差曲线（代表值）

field-8 curve 给出按 input magnitude 划分的精度。选取的 jellyfish（= df = pf）fast-path（variant 0）转折点，按 input threshold：

```text
divide:      input>=0         rel=1.0       ulp=256    (subnormal: 100% error)
             input>=2.35e-38  rel=1.53e-05  ulp=128
             input>=4.70e-38  rel=1.19e-07  ulp=1      (normal range: near-exact)
exponential: input>=0         rel=1.0       ulp=inf
             input>=8.97e-44  rel=1.0       ulp=128
             input>=2.35e-38  rel=1.53e-05  ulp=128    (caps at ~bf16 precision)
logistic:    input>=0         rel=4.0       ulp=6.71e+07
             input>=9.54e-07  rel=2.0       ulp=3.36e+07
             input>=3.82e-06  rel=1.0       ulp=1.68e+07
             input>=3.05e-05  rel=4.88e-04  ulp=8192
tan:         input>=0..32 over 7 segments: rel 4.77e-07→5.96e-08, ulp 8→1

Ghostlite（最佳 generation）显著收紧了这些值，例如 divide input>=2.35e-38 → rel=1.19e-07 ulp=1（jf 上为 1.53e-05/128）；log input>=1.53e-05 → rel=2.38e-07 ulp=1（jf 上为 4.88e-04/4096）；logistic input>=1.19e-07 → rel=1.19e-07 ulp=1（jf 上为 4.0/6.7e7）。

---

Consumer Chain — 精度决策

select_accuracy.cc 中的四个函数把表转换为实现选择：

1. GetTargetAccuracy(HloOpcode, Target&, bool variant) @ 0x1d60dea0 — loader。读取 Target+0x398 = tpu_version()，通过 TpuVersionToString 构建 embed:// URI，调用 tsl::ReadBinaryProto 解析 blob，然后调用 GetAccuracyEntry。返回 StatusOr<AccuracyEntry>；blob 读取失败会在 select_accuracy.cc:143 通过 CreateStatusAndConditionallyLog 产生一个 Status（6acc60406 走的 table-missing 路径）。
2. GetAccuracyEntry(string_view name, AccuracyTable&, bool variant) @ 0x1d60e040 — 线性扫描 table.entries，同时匹配 (entry.name == name)（通过 bcmp）以及 entry+80 处的 variant byte（== variant）。若没有匹配，会构造一个带有 "No accuracy entry found for opcode: "（select_accuracy.cc:129）的 error Status（MakeErrorImpl<5>）。
3. CompareAccuracy(ResultAccuracy&, StatusOr<Entry> fast, StatusOr<Entry> accurate, HloOpcode) @ 0x1d60d7a0 — tolerance check。读取 ResultAccuracy.mode（a2+28）；当 mode 为 TOLERANCE 时，加载 tolerance struct（a2+16，字段在 +0x18/+0x20），并用 vucomisd 与 entry 的顶层 bound（entry+0x40）比较；如果 fast（variant-0）entry 通过则 short-circuit 到它，否则尝试 accurate（variant-1）entry。如果两个 status 都不是 OK，则发出 "Only accurate implementation is available."（:33）或 "Only EUP implementation is available."（:35）；如果没有任何 variant 满足请求，则返回 error Status（MakeErrorImpl<5>，:44），内容为 "No implementation found for requested accuracy: … for opcode: <op>"。
4. GetImplType(Target&, ResultAccuracy&, HloOpcode) @ 0x1d60db80 — master decision。加载两个 variant（GetTargetAccuracy(...,0) 和 (...,1)），并按 ResultAccuracy.mode（a3+28）分支：DEFAULT 返回可用 impl（如果 preferred 缺失则 warning）；HIGHEST 强制 EUP 路径；TOLERANCE（mode == 2）调用 CompareAccuracy。未处理的 mode 不会 abort，而是返回一个 error Status（MakeErrorImpl<3>，:79），内容为 "Unsupported result accuracy for opcode <op>: <message>"。

missing-variant fallback inside GetImplType (also the whole-table-missing case
for 6acc60406, whose embed:// lookup returns NOT_FOUND for both variants):
  EUP entry missing      → WARN "No EUP implementation is available, returning accurate implementation."   (select_accuracy.cc:69)
  accurate entry missing → WARN "No accurate implementation is available, returning default implementation." (select_accuracy.cc:61)
```text

> **NOTE —** `CompareAccuracy`（`0x1d60d7a0`）会把 `ResultAccuracy` tolerance 字段与 entry 顶层 bound（`entry+0x40`）进行 **直接 `vucomisd` 比较**；四个 consumer body 中没有 `ulps × 2^-23` relative-error conversion，也没有乘常量的 `mulsd`。`GetImplType` 对未处理 mode 返回 **error `Status`**，而不是 fatal abort。decompile-confirmed decision logic（CONFIRMED）：两次 `GetTargetAccuracy` 调用；DEFAULT / HIGHEST / TOLERANCE（`mode == 2`）dispatch；`:69`/`:61` 处的两个 missing-variant warning；以及 `:79` 处的 `"Unsupported result accuracy"` error return。

`optional<ResultAccuracy>` 来自 HLO instruction 的 `result_accuracy` attribute，并向下穿过 `LloRegionBuilder` transcendental emitter（例如 `Sln` @ `0x1d534520` 计算 `log(x) = Vlog2(F32, x, accuracy) * ln2`）；在这里，`GetImplType` 会在 EUP 与 `*NoEupF32` software expansion 之间选择；`DecomposeEupInstruction` @ `0x126a0340` 会 lowering 选中的 EUP op。

---

## Accuracy vs Auto-Cast（不同层）

accuracy table 选择的是 *transcendental approximation algorithm*（fast XLU/EUP vs accurate software polynomial）；它从不改变 storage dtype（`type` 始终是 F32）。Storage-dtype 和 matmul-accumulation precision 属于另一层：`AutoMixedPrecision`、`PrecisionConfig`、`MayIncreaseBF16AllReduceAccumulationAccuracy` @ `0x127a22c0`。二者都会馈入 cost model：通过 `CompareAccuracy` 的 fast approximation 会把 op 保持在廉价 XLU 路径，这也是逐代 accuracy 与逐代 transcendental cycle cost 一起改善的原因（sin/cos 在 jf→viperfish→ghostlite 上为 198→154→142，tan 为 219→170→151）。`compiled_with_accuracy_setting` flag @ `0x22581ff0`（`FEATURE_KIND_RESULTS_ACCURACY_SETTING`）记录 executable 是否带 explicit accuracy setting 编译，用于 cache-key/replay。

---

## 相关组件

| 名称 | 关系 |
|---|---|
| `tpu::TpuVersionToString` @ `0x20b3a480` | 为 `embed://` URI 解析 codename |
| `tsl::ReadBinaryProto` @ `0x20d02060` | 解析嵌入式 blob（每次调用都会重新解析） |
| `LloRegionBuilder::Sln` / `Vlog2` / `*NoEupF32` | `GetImplType` 在其间路由的 emitter |
| `DecomposeEupInstruction` @ `0x126a0340` | lowering 选中的 EUP transcendental |

## 交叉引用

- [Codename Matrix](tpu-version-codename-matrix.md) — version stamp 依赖的 `tpu::TpuVersion` ↔ codename 映射
- [Per-Codename Constants](per-codename-hw-constants.md) — 逐代 hardware-constant 表（兄弟逐代数据源）
- [EUP Correction Coefficients](../cost/eup-correction-coeffs.md) — accurate variant 调用的 EUP-path correction term