Error/Status 字符串模板

本页所有偏移、符号和计数均适用于 libtpu-0.0.40-cp314 wheel 中的 libtpu.so（libtpu_lts_20260413_b_RC00，ELF x86-64，781,691,048 字节，BuildID md5 89edbbe81c5b328a958fe628a9f2207d，未 stripped）。其他构建会有所不同。

摘要

libtpu 的诊断信息由内嵌在 .rodata 中的格式字符串模板构成，并在调用点填充。有两种格式化习惯用法和一种包装约定。主导习惯用法是 C printf 风格，即 absl::StrFormat("…%s…%d…", args) 和 LOG(...) << absl::StrFormat(...)，覆盖了绝大多数表面；少数习惯用法是 absl 位置替换，即 absl::Substitute("…$0…$1…", args)，主要限于 megascale DCN runtime、collective 缓冲区大小校验器，以及一个静态链接的 protobuf descriptor 校验器。随后，这些模板大多会由三个 factory 家族之一包装成 absl::Status（见 Status-Code 映射），因此同一个字符串同时是人类可读错误，也是经由 PJRT 传回 JAX/XLA 用户的 StatusOr<T> 载荷。

本页是参考目录，不是算法轨迹。它的价值在于按组去重的真实模板表，并解释其占位符、在字节已确认时标注其 absl::Status code。恢复出的表面大约有 2,937 个不同的 error/status 模板；本页不会全部复现，而是记录这个空间的结构（占位符语法、子系统划分、三种 Status 构造习惯用法、字节已确认的参数类型），并为每个子系统给出有来源锚点的代表性样本，使重新实现者可以预测未列出模板的形态。

占位符语义是贯穿全篇的主线。%s 几乎从来不是原始 C 字符串：它是某个对象的 ToString() / AbslStringify，例如 Shape、Layout、HloInstruction 名称、设备名称或 opcode 助记符。%d/%lld/%zu 是维度、ordinal、计数或字节大小。指针（%p）和浮点数（%f）只出现在低层 driver 和 cost-model 路径中。本目录会逐个模板指出每个占位符的含义，因为单看 printf 类型（%s = char*）会隐藏真实的 C++ 参数。

注意 — 此处的 error/status 模板不同于面向用户的提示字符串（以 "try…"、"consider…" 这类建议形式表达，而非失败），后者位于 hint-strings.md；也不同于 internal-pass-names.md 上的内部 pass-name 字符串（pipeline 阶段标识符，不是诊断）。本页负责 error/status 格式字符串目录、占位符语义和 status-code 映射。


不同的 error/status 模板	~2,937（~2,799 个 printf 风格，~138 个位置 `$N`/`%v`）
Status 构造习惯用法	3（`<Code>StrCat` factory、`MakeErrorStream`、散文式 `absl::<Code>Error`）
字节已确认的参数类型 factory	99 个 `xla::<Code>StrCat<Types…>` 实例化
`MakeErrorStream` 返回类型操作	390 个 `MakeErrorStreamWithOutput<T>`（每个 `StatusOr<T>` 一个）
最常见占位符	`%s`（~2778）、`%d`（~2031）、`%u`（116）、`%zu`（84）、`%x`（55）、`%lld`（52）
最大子系统块	Compile / HLO / verifier（~875 个候选）
经反编译抽查确认	12/12 个代表性模板（见报告）

如何阅读本目录

两种格式化习惯用法

text

printf-style   "Argument to Cholesky must have rank >= 2; shape was %s"
               filled by  absl::StrFormat(template, shape.ToString())
               or         LOG(ERROR) << absl::StrFormat(template, …)

positional     "ALL_TO_ALL not supported when buffer size is not divisble
                by number of endpoints. Buffer Size: $0 Number of
                endpoints: $1. MegascaleInfo: $2"
               filled by  absl::Substitute(template, size, n, info)

printf 表面使用 % spec；位置表面使用 $0、$1、…（以及少量 %v，即 absl 的 stringify-any 标记）。两者绝不会混在同一个模板里。只依据 % spec 的重新实现者会完全漏掉约 138 个位置模板。

占位符语法

所有模板中 % spec 出现次数的完整分布明显偏向字符串和整数：

Spec	计数	C 类型	真实 C++ 参数（典型）
`%s`	~2778	`char*`	`ToString()` / `AbslStringify` 结果，例如 Shape、Layout、instruction 名称、device 名称、opcode
`%d`	~2031	`int`	维度索引、ordinal、计数、status/state enum
`%u`	116	`unsigned`	不可能为负的计数或索引
`%zu` / `%zd`	84 / 19	`size_t` / `ssize_t`	字节大小或元素计数
`%x` / `%#x` / `%02x`	55 / 19 / 11	`unsigned`（十六进制）	地址片段、寄存器/位掩码、chip ID 字节
`%lld` / `%ld` / `%lu` / `%llu`	52 / 48 / 40 / 4	`long long` / `long` / `unsigned long`	大字节大小或 64 位计数
`%p`	32	`void*`	缓冲区/driver 对象指针（仅 driver + cost-model 路径）
`%c`	24	`char`	大括号/方括号字面量或轴字母（`'{'`、`'}'`、`'X'`）
`%f`	19	`double`/`float`	比率/分数（cost model、hbm-fraction 配置）
`%v`	6	absl-stringify	带有 `AbslStringify` overload 的对象（例如 `TensorCoreBundle`）

陷阱 — %s 并不说明程序中的参数是字符串。它只说明调用点传入了 const char*，而这几乎总是某个对象 ToString() 的产物。重新实现时，%s 背后的真实类型只能在调用点恢复（例如 "%s vs. %s" 的操作数到底是 Shape::ToString() 还是 HloInstruction::name()）；本目录会注明可能类型，但不会对 printf 家族做字节确认。

置信度与归属

标记为 CERTAIN 的模板已在反编译调用点逐字抽查确认。子系统分组是对文本的关键词分类，不是逐模板调用点轨迹；对文本内容为 HIGH 置信，对精确所属 pass 为 MEDIUM。status code 列只有对 99 个 <Code>StrCat factory（mangled symbol 命名 code）和 390 个 MakeErrorStream 位置（macro 命名 code）为 CERTAIN；其他所有模板的 code 都依据文本和所属子系统的主导习惯用法推断，并相应标记为 MEDIUM 或 LOW。

Compile / HLO / Codegen / Verifier — ~875 个模板

最大的单一块。这些是 JAX/XLA 用户在程序格式错误时遇到的 shape-inference 和 HLO-verifier 诊断：rank 检查、维度范围检查、shape 相等检查、layout 不匹配。几乎全部重度使用 %s/%d，并且几乎全部包装 InvalidArgument（少数 RET_CHECK → Internal）。verifier 的特征是 vs. 比较风格，即 "%d vs. %d"、"%s vs. %s"，会出现在两个量必须匹配的任何位置。

偏移	模板（`%` spec）	占位符	Code
`0x857d595`	`Argument to Cholesky must have rank >= 2; shape was %s`	`%s`=Shape	InvalidArgument
`0x857c6ef`	`Argument to symmetrize must have >= 2 dimensions, got %s`	`%s`=Shape	InvalidArgument
`0x8580369`	`All reduced tensors must have the same dimension. Tensor 0 has shape %s, Tensor %d has shape %s`	`%s`=Shape, `%d`=索引, `%s`=Shape	InvalidArgument
`0x85803c9`	`All operands to AfterAll must be tokens; operand %d has shape %s`	`%d`=operand idx, `%s`=Shape	InvalidArgument
`0x872a259`	`broadcast_dimensions contains invalid value %d for result with rank %d`	`%d`=值, `%d`=rank	InvalidArgument
`0xa02f654`	`Broadcast dimension %d mismatch: %d != %d; %s and %s.`	`%d`=dim, `%d`/`%d`=大小, `%s`/`%s`=Shapes	InvalidArgument
`0xa02c26f`	`Cannot concatenate arrays that differ in dimensions other than the one being concatenated. Dimension %d in both shapes must be equal (or compatible): %s vs %s.`	`%d`=dim, `%s`/`%s`=Shapes	InvalidArgument
`0xa02f9ba`	`Cannot bitcast types with undivisible bit-widths: %s => %s.`	`%s`/`%s`=PrimitiveType	InvalidArgument
`0xa01cdc3`	`Bitcast requires a new on-device shape to have the same size of %d bytes, but got %d bytes.`	`%d`/`%d`=字节大小	InvalidArgument
`0xa030ae8`	`Cannot infer shape: attempting to index into non-tuple: %s.`	`%s`=Shape	InvalidArgument
`0x857b3fa`	`sharding's tile count and device count does not match: %d vs. %d; shape=%s, sharding=%s`	`%d`/`%d`=计数, `%s`/`%s`=Shape/Sharding	InvalidArgument
`0x858b317`	`Arguments to TriangularSolve have shapes with different ranks: %s vs. %s`	`%s`/`%s`=Shapes	InvalidArgument
`0xa0a2a82`	`Binary op shape inference: %s; lhs: %s; rhs: %s is not implemented.`	`%s`=op, `%s`/`%s`=Shapes	Unimplemented
`0x8728000`	`Binary op expects 2 operands, but got %d`	`%d`=计数	RET_CHECK→Internal
`0x858400c`	`Bad scalar opcode in slot 0, opcode: %d bundle: %v, bits: %s`	`%d`=opcode, `%v`=TensorCoreBundle, `%s`=bits	InvalidArgument
`0x857b280`	`Cannot feed constants into bundle packer. Copy them to registers first. instr=%s`	`%s`=instruction	Internal
`0x8584e00`	`Cannot find a free bundle slot in bundle %s: %s`	`%s`/`%s`=bundle/reason	Internal

注意 — 0x858400c（"Bad scalar opcode in slot 0…"）在反编译中解析到 platforms_deepsea::jellyfish::isa::DecoderBcsDf::DecodeScalar0Slot（以及一个逐代 DecoderJf twin）。它的 %v 是 TensorCoreBundle 的 AbslStringify，也是 6 个带 TensorCoreBundle 的 <Code>StrCat factory 之一（见参数类型解码）。逐代 decoder 变体（gxc/gfc、gxc/glc、vxc）各自发出自己的副本，因此文本几乎相同的模板、即使 prose 相同，也是在 rodata 中分离的字符串。

Scheduler / Fuel / FIFO — ~15 个模板

Latency-hiding scheduler、annotation 范围检查、--xla_fuel 预算，以及设备端 FIFO push/pop 顺序。annotation 模板依赖 %c 表示大括号/方括号字面字符；FIFO 模板串联 %s :: %s 来附加 instruction 上下文。

偏移	模板	占位符	Code
`0x8796ba8`	`annotation arg must be in correct order as given; expected %c{%d%c but got %c{%d%c`	`%c`=brace, `%d`=id	Internal
`0x8571b33`	`annotation %c{%d%c is out of bounds`	`%c`=brace, `%d`=id	Internal
`0x858a654`	`annotation range was not closed; expected %c}%c: %s`	`%c`=brace, `%s`=context	Internal
`0x857fa91`	`async-done for %s must be scheduled before %s`	`%s`/`%s`=instructions	RET_CHECK→Internal
`0x857fabf`	`async-done for %s must be scheduled on core %d before %s`	`%s`=instr, `%d`=core	RET_CHECK→Internal
`0x858a9b8`	`Cannot schedule FIFO pop instruction when the FIFO is empty %s :: %s`	`%s :: %s`=instr context	Internal
`0x857ad7e`	`Cannot schedule FIFO push instruction when the FIFO is full. FIFO name: %s. (element count %d vs %d). %s :: %s%s`	`%s`=name, `%d vs %d`=计数	Internal
`0xa02bf0c`	`Conflicting schedule type requirements in computation rooted at %s.`	`%s`=computation	Internal
`0xa086733`	`Reference instruction %s was not found in the schedule.`	`%s`=instruction	Internal
`0x862868f`	`GVN: Not replacing %s because GVN is out of fuel`	`%s`=instruction	（LOG，非 Status）
`0x8628660`	`halt before %s because lowering is out of fuel`	`%s`=instruction	（LOG，非 Status）
`0xa03ca6a`	`Illegal value for --xla_fuel. Saw %s, but expected token %s to be an integer.`	`%s`/`%s`=value/token	InvalidArgument

注意 — 0xa03ca6a（"Illegal value for --xla_fuel…"）已确认位于 xla::MakeDebugOptionsFlags flag 解析闭包内，它是 flag 值校验器，也就是 flag-name 侧所记录 --xla_fuel flag 的错误对应项。0x8628660/0x862868f 两个 "out of fuel" 字符串是诊断 LOG 输出，不是 Status 载荷（没有 factory 包装）；这里为完整性收录，但不应假定它们能恢复为 Status code。

MSA / Memory / Allocation — ~79 个模板

Memory-space assignment、prefetch/alternate-memory、HBM defragmentation 和 heap allocator。字节大小占位符占主导（%lld、%zu）。两个超预算模板包装 ResourceExhausted；不匹配/校验模板包装 Internal。

偏移	模板	占位符	Code
`0xa030cd9`	`AllocateBufferForMemorySpace: Unsupported memory space: %s.`	`%s`=memory space	InvalidArgument
`0x858aa58`	`Allocation (size=%lld) would exceed memory (size=%lld) :: %s :: %s`	`%lld`/`%lld`=大小, `%s :: %s`=context	ResourceExhausted
`0xa083c62`	`BufferAllocation::Slice for instruction %s at index %s cannot be determined at compile-time.`	`%s`/`%s`=instr/index	Internal
`0x857ce50`	`DefineBuffer: Mismatch in memory spaces: %s vs %s`	`%s vs %s`=spaces	Internal
`0x8584ecd`	`Error defragmenting HBM %s: %s`	`%s`/`%s`=region/reason	Internal
`0xa1300d0`	`Failed to allocate %zu bytes. Memory limit: %zu bytes. Used: %zu bytes.)`	`%zu`×3=req/limit/used	ResourceExhausted
`0x8728f50`	`Invalid HBM offset %d`	`%d`=offset	InvalidArgument
`0x872cb26`	`Invalid memory space for input memory space colors: %d`	`%d`=color	InvalidArgument
`0xa01cf08`	`Out of memory allocating %d bytes.`	`%d`=字节大小	ResourceExhausted
`0xa09a63e`	`Number of bytes %lld allocated must be a multiple of chunk size %lld.`	`%lld`/`%lld`=size/chunk	InvalidArgument
`0x857eed1`	`Register allocator verification failure: live range %s; instruction %s`	`%s`/`%s`=range/instr	Internal
`0xa02b12f`	`Scoped allocation with size %s and limit %s exceeded scoped %s limit by %s.`	`%s`×4=sizes/labels	ResourceExhausted

注意 — 不要把近似相同的模板 0xa13e8e5（"Failed to allocate node (%zu bytes). Memory limit: %zu [bytes]. Used: %zu [bytes].)"）归因于 TPU memory-space assignment。它位于 perfetto::protovm::RwProtoCursor::CreateNodeFromField，即 Perfetto tracing library 的 arena allocator，而非 TPU MSA 路径。真正的 TPU heap-allocator 超预算字符串是 0xa1300d0（"Failed to allocate %zu bytes…"，上表行），已在 xla::AlignedAllocator::Allocate 中确认。

ICI / Collective — ~90 个模板

Inter-chip-interconnect 链路健康、routing、GTC 同步，以及 collective（all-reduce / all-gather / reduce-scatter / all-to-all）缓冲区大小校验器。本块混合了 printf 风格的 ICI driver 错误和位置 $N collective 校验器（带 megascale 标记，见下一节）。Code 倾向于 Internal 和 DeadlineExceeded。

偏移	模板	习惯用法	占位符	Code
`0xa0b3abc`	`Cannot find unicast link next hop routing table for link port %d.`	printf	`%d`=port	Internal
`0xa030c51`	`Coordinate assignment failed for the slice's target %s ICI network because there are chips disconnected from the rest of the slice: %s.`	printf	`%s`/`%s`=target/chips	Internal
`0xa0d5412`	`Detected ICI link failures along %d dimensions, but only 1-dimensional link fault is allowed..`	printf	`%d`=dim 计数	FailedPrecondition
`0xa05e59a`	`Failed to add link information: chip %d already has a %c direction link.`	printf	`%d`=chip, `%c`=axis	Internal
`0x855f8c6`	`Failed to detect GTC reset before timeout %s expires`	printf	`%s`=duration	DeadlineExceeded
`0x8727005`	`Failed to turn down ICI link %d during slice reset, state=%d`	printf	`%d`=link, `%d`=state	Internal
`0x871106b`	`GTC failed to converge (max diff %d > %d) before timeout (%s) expired`	printf	`%d > %d`=diff, `%s`=timeout	DeadlineExceeded
`0x872a345`	`Hop ID %d is out of bound of ICI route path with length %d`	printf	`%d`/`%d`=hop/len	Internal
`0x8583d20`	`ICI Probe failed. local port: %d name: %s took %d us. status: %s`	printf	`%d`=port, `%s`=name, `%d`=us, `%s`=status	Internal
`0xa0a81a9`	`ICI resiliency only allow 1-dimensional link failures, but link failures along %d dimensions are discovered.`	printf	`%d`=dim 计数	FailedPrecondition
`0xa0ba533`	`ICI routing failed to retrieve %dth hop dimension from bit encoded cache data.`	printf	`%d`=hop	Internal
`0x9a573e9`	`ALL_REDUCE Output buffer size is not == Input buffer size. Input size: $0 Output size: $1 Group Size $2 Key: $3 Module: $4 MegascaleInfo: $5`	positional	`$0..$5`=sizes/key/module/info	InvalidArgument
`0x9c142f9`	`ALL_GATHER Input buffer size is not (Output buffer size / group size). Input size: $0 …`	positional	`$0…`=sizes/info	InvalidArgument
`0x9c14380`	`REDUCE_SCATTER Output buffer size is not (Input buffer size / group size). …`	positional	`$0…`=sizes/info	InvalidArgument
`0x9d1493e`	`ALL_TO_ALL not supported when buffer size is not divisble by number of endpoints. Buffer Size: $0 Number of endpoints: $1. MegascaleInfo: $2`	positional	`$0`/`$1`/`$2`	Unimplemented

注意 — 0x8583d20（"ICI Probe failed…"）已在 asic_sw::driver::deepsea::ici::SliceConfiguration::GetLocalTopology 中确认；0x9d1493e（"ALL_TO_ALL not supported…"）已在 xla::megascale::runtime::HostCommandSchedulerFactory…GenerateCommunicationIrsFromTransferRegistry 中确认。collective 缓冲区大小校验器虽然位于 ICI/collective 路径上，却是位置形式（$N）：它们由 megascale runtime 发出，而这里是位置习惯用法的大本营。

Megascale（DCN Runtime / Aggregator）— ~21 个模板

跨主机 data-center-network 协调：barrier 参与者记账、损坏缓冲区检测器、launch-id timeout，以及 coordinator 的错误摘要。这是位置 $N 习惯用法的核心。coordinator 的 hang 摘要会为每个 cause branch 发出一种 prose 变体。

偏移	模板	占位符	Code
`0x9e6f85e`	`Extra barrier participant. Expected: $0 Message $1`	`$0`=expected, `$1`=msg	Internal
`0x9e6fa4d`	`Mismatched number of barrier participants: Expected: $0 Msg: $1`	`$0`=expected, `$1`=msg	Internal
`0x9d149cb`	`MegaScale Corrupted Buffer Detected. Key: $0 Checksum at Sender: $1 Current checksum: $2`	`$0`=key, `$1`/`$2`=checksums	DataLoss/Internal
`0x9b273a4`	`Timed out waiting for $0 graphs to complete at launch_id $1. Already completed: $2. StepGloballyInProgress: $3 Timeout: $4`	`$0..$4`=count/id/state	DeadlineExceeded
`0xa122d03`	`MegaScale devices cannot be queried except from jax. (%d)`	`%d`=error code	FailedPrecondition

coordinator 的 hang 摘要会针对每个 cause branch（BAD_TPU_CHIP、BAD_SC_CHIP、DATA_INPUT_STALL、DIFFERENT_MODULE、FINGERPRINT_MISMATCH、NETWORKING_ISSUE、PROGRAM_NOT_QUEUED、UNKNOWN_CAUSE）发出 prose "Megascale detects a hang that is likely caused by …"，以及面向 operator、可行动的后续建议（"Please remove the hosts from the fleet and restart the workload"、"Please check the workers to make sure the data input pipeline is working properly"）。abort 路径见 Fatal / Abort 表面。

SparseCore / Embedding — ~68 个模板

SparseCore（xla_sc_）和 BarnaCore embedding 配置：alignment 要求、table/feature 计数、partitioner objective enum，以及 SMEM row-pointer 预算。面向用户（模型配置错误），且几乎全部为 InvalidArgument。

偏移	模板	占位符	Code
`0xa0a9715`	`barna_core_infeed_queue_hbm_address must be %d-byte aligned.`	`%d`=alignment	InvalidArgument
`0xa0b6320`	`barna_core_infeed_queue_hbm_size must be a multiple of %d.`	`%d`=multiple	InvalidArgument
`0xa030dd4`	`Could not find valid TPU batch of length at least %d at position %d for row %d. The embedding work in one sample exceeds what the BarnaCore can process: %s. %s.`	`%d`×3=len/pos/row, `%s`/`%s`=detail	InvalidArgument
`0x872d950`	`Dynamic learning rate tag: %d not found in the TPU embedding configuration, instead found: %d. tag set size: %d`	`%d`×3=tag/found/size	InvalidArgument
`0xa02cd93`	`Embedding table is expected to have element type %s or %s.`	`%s`/`%s`=types	InvalidArgument
`0x86fa1ad`	`Failed to parse TPU embedding partitioner optimization objective "%s". Valid options: performance, hbm_usage, hybrid`	`%s`=value	InvalidArgument
`0xa11773a`	`hbm_limits_for_embeddings.min_fraction (%f) must be <= hbm_limits_for_embeddings.max_fraction (%f)`	`%f`/`%f`=fractions	InvalidArgument
`0xa0d36d1`	`Invalid num_features: %d found for table: %s in the TPU embedding configuration. Valid values are >0.`	`%d`=计数, `%s`=table	InvalidArgument
`0xa0b7076`	`Logical replicas must evenly divide the SparseCores in the system. logical_replicas = %d, physical_sparse_cores = %d.`	`%d`/`%d`=计数	InvalidArgument
`0xa069f4e`	`Number of TPU tables on row: %d exceeds what the BarnaCore hardware supports: %d > %d. This is mostly likely a result of incorrect partitioning.`	`%d`×3=row/got/max	InvalidArgument
`0xa0ff40e`	`Row pointers would exceed available SCS Smem (%d bytes > %d bytes)`	`%d`/`%d`=used/avail	ResourceExhausted
`0xa07d172`	`Scatter operand has %d elements, which exceeds the 32-bit limit. Unsupported on SparseCore.`	`%d`=计数	Unimplemented

注意 — 0xa0d36d1（"Invalid num_features…"）已在 tensorflow::PopulateMissingFieldsInTPUEmbeddingConfig 中确认。0xa11773a 中的 %f 浮点数是少见情形：%f 真的是配置比率，而非 cost-model 内部值。

Runtime / Driver / PJRT — ~177 个模板

driver state machine、device/ordinal 校验、firmware-queue 转换、DMA-buffer 记账，以及 PJRT C-API 边界。习惯用法最混杂的一块：这里会出现 %p 和 errno（%d），code 在 FailedPrecondition（state-machine guard）和 Internal 之间分裂。

偏移	模板	占位符	Code
`0xa0b7366`	`Attempted to register programmable interrupt with bad index: %d. Number of programmable interrupts: %d.`	`%d`/`%d`=index/count	InvalidArgument
`0xa0430a3`	`Cannot remove a driver for %s, was not found in map.`	`%s`=driver name	NotFound
`0xa077a78`	`Cannot transition to %s: the firmware queues are not in %s state; they are in %s state.`	`%s`×3=states	FailedPrecondition
`0x96c33b2`	`Can't close driver while in state %s; are multiple threads trying to open / close?`	`%s`=state	FailedPrecondition
`0x94b68ce`	`Can't get the optimized program for executable \`%s`: MPMD execution is not supported by PJRT C API`	`%s`=executable	Unimplemented
`0xa09fdcd`	`Chip count (%d) is not supported.`	`%d`=计数	InvalidArgument
`0x872d1a0`	`Close of core dump fd failed with errno: %d`	`%d`=errno	Internal
`0xa0a9937`	`%d DMA buffers were still outstanding when the driver was re-opened. These buffers must be unmapped before the driver can be re-opened.`	`%d`=计数	FailedPrecondition
`0xa0d10bd`	`Device id '$0' is out of bound. Number of devices is $1.`	`$0`/`$1`=id/count	InvalidArgument
`0x8679159`	`device ordinal value (%d) must be non-negative`	`%d`=ordinal	InvalidArgument
`0xa1a96ba`	`executable is built for device %s of type "%s"; cannot run it on device %s of type "%s"`	`%s`×4=device/type	InvalidArgument
`0xa00ab4b`	`Expected %d chips per tray, actually found a tray with %d chips.`	`%d`/`%d`=expected/found	FailedPrecondition
`0x858a3de`	`failed initializing StreamExecutor for device ordinal %d: %s`	`%d`=ordinal, `%s`=reason	Internal
`0xa09b555`	`Failed to convert multipod chip id %d to single-pod chip id.`	`%d`=chip id	Internal

注意 — 0x8679159（"device ordinal value (%d) must be non-negative"）已在 stream_executor::StreamExecutorAddressAllocator::GetStreamExecutor 中确认。许多 runtime 模板会进入 executor 的 async stream 路径；参见 execute-async-on-stream.md，了解它们在入队期间出现的位置。

PJRT-C-API / protobuf-descriptor（链接库，非 TPU）

位置 $N 家族也由一个静态链接的 protobuf extension-declaration 校验器填充。这里为完整性收录，但它们不是 TPU code，不应归因于 libtpu 自身表面：

偏移	模板	说明
`0xa0cf40b`	`"$0" extension field $1 is expected to be $2.`	protobuf descriptor 校验器
`0xa0eba97`	`"$0" extension field $1 is expected to be type "$2", not "$3".`	protobuf descriptor 校验器
`0xa0d0fab`	`$0 cannot declare both \`metadata` and `declaration` as extension declaration for extension #$1.`	protobuf descriptor 校验器

CHECK / RET_CHECK / 自检模板

absl CHECK/QCHECK/DCHECK 家族和 XLA 的 TPU_RET_CHECK 不携带完整消息模板：它们发出固定前缀，随后追加字符串化的源表达式和任何流式传入的 << "msg"。比较 macro（CHECK_EQ/NE/GE）会用 verifier 块中反复出现的 "%d vs. %d" / "%s vs. %s" 格式追加操作数值。

偏移	模板	Macro / 来源
`0xa1a64de`	`Check failed: '`	absl `CHECK`/`QCHECK` 前缀（带引号）
`0xa1f4a87`	`Check failed in`	带 file/line 的 absl CHECK
`0xa285fb1`	`Check failed:`	absl CHECK 前缀（无引号）
`0xa183292`	`TPU_RET_CHECK failure (`	XLA TPU RET_CHECK macro
`0xa0ab3ca`	`Hostname Verification Check failed.`	gRPC TLS hostname-verify CHECK
`0xa2300c5`	`MakeErrorStream destructed without getting absl::Status:`	XLA status_macros 自检
`0xa2300ff`	`MakeErrorStream shift called after getting absl::Status:`	XLA status_macros 自检
`0xa27f1b3`	`MakeErrorStream got absl::Status more than once:`	XLA status_macros 自检

陷阱 — CHECK(expr) << "msg" macro 会把源表达式文本（spmem_buffer_type != nullptr、dynamic_size, nullptr）内联进 .rodata。大量 C++ 源码字面片段，甚至整段 lambda body，都会作为副作用出现在字符串表中。这些是 CHECK 条件证据，不是 printf 意义上的错误模板，已从模板计数中排除。扫描字符串寻找 "templates" 的重新实现者必须过滤它们，否则会被源码片段淹没。

0xa183292 已在 tpu::internal::RetCheckFailSlowPath 中确认；0xa2300c5 已在 xla::status_macros::MakeErrorStream::Impl::~Impl 中确认（当一个已构造但未消费的 Status 被丢弃时触发的析构自检）。

Fatal / Abort 表面 {#fatal--abort-surface}

有意 abort 的表面很小且集中：ICI 硬故障、megascale coordinator abort、内部 bug LOG(FATAL)，以及一个非 TPU 专属的 library guard。

偏移	模板	路径
`0xa1e7cb3`	`!!!! FATAL ERROR !!!! for`	ICI `FatalErrorCheck`
`0x8864055`	`!!!! FATAL ERROR !!!! observed errors are: [`	`AsyncDriver::HandleFatalError` composite
`0xa046045`	`Fatal error occurred. Data links will go down.`	ICI hard-failure marker
`0xa1b3810`	`FATAL ERROR RECEIVED FROM HARDWARE!!!`	hardware fatal interrupt
`0x8a2941d`	`Fatal error in creation of RWB Fusion. Please file a bug with XLA-TPU`	fusion internal-bug `LOG(FATAL)`
`0xbe7d460`	`FATAL ERROR: This binary was compiled with <isa> enabled, but this feature is not available on this processor (go/sigill-fail-fast).`	absl CPU-feature startup guard（12 个 ISA 变体）

megascale coordinator 的 abort prose："Aborting the coordinator after collecting errors from all workers as megascale_error_reporter_abort_on_hang is set to true. All workers will also abort after they detect the coordinator is shutdown."，是由 flag gated 的 LOG(FATAL)。

注意 — 0xbe7d460 家族（12 个变体：aes、avx、mmx、pclmul、popcnt、sse、sse2、sse3、sse4.1、sse4.2、ssse3、…）是 absl CPU-feature startup guard，即 library abort：如果 host CPU 缺少该 binary 构建时启用的 ISA，它会在任何 TPU code 运行前触发。它不是 TPU 诊断。0x8a2941d（"Fatal error in creation of RWB Fusion…"）已在 xla::jellyfish::TpuInstructionFusion::RwbFusionHelper 中确认，是带有 "file a bug" 标记的真正内部 bug fatal。

Status-Code 映射 {#status-code-mapping}

模板通过三种 factory 习惯用法之一变为 absl::Status。只有前两种在 binary 中命名 code；prose 家族和裸 printf-into-LOG 路径需要调用点反汇编才能确认 code。

习惯用法	计数	含义	Code 归属
`xla::status_macros::MakeErrorStream`	390	`RET_CHECK(c) << …` / `return InvalidArgument(…) << …`；390 个 `MakeErrorStreamWithOutput<T>` conversion op，每个 `StatusOr<T>` 返回类型一个。按调用点计数是主导 XLA 习惯用法。	code 由该位置的 macro 命名
`xla::InvalidArgumentStrCat<…>`	38	交错 literal/typed-arg factory	CERTAIN → InvalidArgument
`xla::UnimplementedStrCat<…>`	31	同上	CERTAIN → Unimplemented
`xla::InternalStrCat<…>`	29	同上	CERTAIN → Internal
`xla::ResourceExhaustedStrCat<…>`	1	同上	CERTAIN → ResourceExhausted
`absl::<Code>Error("…")` prose	1+	直接 prose factory（例如 `absl::InternalError("Invalid error type")`）	code 由函数命名

<Code>StrCat factory 是唯一一个无需反汇编即可恢复 format-arg C++ 类型的地方：每个实例化都会发出 Itanium-mangled symbol _ZN3xla<len><Code>StrCatIJ<typepack>EEC2E…，其中 <typepack> 会逐字节命名每个交错 literal segment 和 typed argument。

参数类型解码（mangled type pack）

text

RA<N>_Kc   const char (&)[N]  — a literal segment of N chars (N counts the
                                 NUL, so the visible literal is N-1 chars)
RA<N>_S1_  another const char (&)[N] — a later literal (S1_ back-references
                                 the already-named const char type)
m / l / i / f          unsigned long(size_t) / long / int / float, by value
Rm / Rl / Ri / Rf      &  of each
RKm / RKi              const unsigned long& / const int&
NSt…basic_string…      std::string by value;  RKNSt… = const std::string&
NSt…basic_string_view… std::string_view
N3tpu10TpuVersionE     tpu::TpuVersion (enum), by value
N…isa16TensorCoreBundleE  the per-gen ISA TensorCoreBundle (rendered via its
                          AbslStringify → the %v sigil)
N4absl8StatusOrI…E     absl::StatusOr<…>, by value

解码示例（每个都是真实的实例化 symbol）：

text

InvalidArgumentStrCatIJRA74_KcmEE
    → InvalidArgument("<73-char literal>", size_t)
InvalidArgumentStrCatIJRA16_KcRfRA20_S1_EE
    → InvalidArgument("<15ch>", float&, "<19ch>")     (the only float-bearing
      factory family — a cost/ratio message)
InvalidArgumentStrCatIJRA19_KcRlRA38_S1_RKiEE
    → InvalidArgument("<18ch>", long&, "<37ch>", const int&)
InvalidArgumentStrCatIJRA26_KcN3tpu10TpuVersionEEE
    → InvalidArgument("<25ch>", tpu::TpuVersion)
InvalidArgumentStrCatIJRA65_Kc…isa16TensorCoreBundleEEE
    → InvalidArgument("<64ch>", const TensorCoreBundle&)   (the "Bad scalar
      opcode … bundle: %v" family — per-gen gxc/gfc, gxc/glc, vxc variants)
UnimplementedStrCatIJRA253_KcEE
    → Unimplemented("<252-char literal>")   (the longest single-literal
      Unimplemented message — a "not supported" explanation block)

99 个 factory 中字节已确认的 arg-type 频率：std::string 29、std::string_view 15、long 13、TensorCoreBundle 6、size_t 5、StatusOr<…> 3、const int& 2、float 2、TpuVersion 1。即使在字节已确认层面，args 也绝大多数是类字符串 + long；float 和 pointer 极少，这与上面的 printf-spec 分布一致。

StatusCode 关键词分布

字符串表关键词出现次数（已过滤噪声），这是表面偏向哪些 code 的文本信号上界，不是逐模板计数：

text

InvalidArgument 300 | Unimplemented 169 | NotFound 75 | FailedPrecondition 56
ResourceExhausted 21 | Unavailable 15 | OutOfRange 14 | Aborted 8
DeadlineExceeded 3 | AlreadyExists 3 | PermissionDenied 2 | DataLoss 1

陷阱 — 原始 Internal 关键词计数（~16,823）主要由内联 __FILE__ 字符串中的 /internal/ source-path 组件贡献，而非错误模板。不要把它解读为 Internal-status 计数。字节已确认的 Internal 总数是 29 个 InternalStrCat factory 加上 MakeErrorStream Internal 位置。

面向用户与内部的划分

prose 本身会标示受众。面向用户（operator / JAX-user 可行动、无 bug id）：HLO verifier shape/rank/dimension 块、--xla_fuel flag 值错误、带有效值提示的 embedding-config 错误、megascale 摘要后续建议，以及 device/executable 不匹配。内部（XLA-bug 标记）：带有 b/<id>、go/<link>、"please file a bug" 或 "should not happen" 的字符串。

偏移	内部 bug 标记模板	标记
`0x858c562`	`Kernel body fingerprint collision detected for key: … Please file a bug with the XLA team …`	"file a bug"
`0x8a2941d`	`Fatal error in creation of RWB Fusion. Please file a bug with XLA-TPU`	"file a bug"
`0x96c1211`	`XLA has not implemented dynamic sized slice with non-trival stride yet. Please file a bug against XLA`	"file a bug"
`0x96c12d4`	`Unimplemented reduce-window in fusion cost modeling. Please file a bug with XLA`	"file a bug"
`0x9fd476d`	`tightened domains are empty. This should not happen except if we proven infeasibility or optimality.`	"should not happen"
`0xa0c9452`	`Encountered unexpected layout … This should not happen - please file a bug against XLA.`	"should not happen" + "file a bug"
`0x99e1405`	`Close() appears to be hanging, this might be a deadlock see b/147787375`	`b/<id>`

注意 — b/<id> 和 go/<link> token 是最强的内部信号。它们也标记已知空缺 TODO（"TODO(b/157237781) support VFIO device %s"、"TODO: b/475913712 - Expected Gather indices to be bitpacked …"），这些不是活动错误，但当其受保护路径被命中时会变为 Unimplemented 消息。

概览：各子系统模板数

子系统	候选	主导习惯用法	主导 code	受众
Compile / HLO / verifier	~875	StrFormat `%s`/`%d` + RET_CHECK	InvalidArgument	面向用户
Runtime / driver / PJRT	~177	StrFormat `%s`/`%d`/`%p` + errno	Internal / FailedPrecondition	混合
ICI / collective	~90	StrFormat `%d`/`%s` + `$N` (msc)	Internal / DeadlineExceeded	混合
MSA / memory / allocation	~79	StrFormat `%zu`/`%lld`	ResourceExhausted / Internal	混合
SparseCore / embedding	~68	StrFormat `%d`/`%s`/`%f`	InvalidArgument	面向用户
Megascale (DCN runtime)	~21	Substitute `$0`/`$1`	Internal / LOG(FATAL)	混合
Scheduler / fuel / FIFO	~15	StrFormat `%s`/`%d`/`%c`	Internal / RET_CHECK	混合
PJRT-C-API / protobuf-desc	(lib)	Substitute `$0`/`$1`	InvalidArgument	（非 TPU）
不同模板总数	~2,937	2,799 printf + 138 positional	—	—

陷阱 — 子系统计数是在 prose 关键词集上按 first-match-wins 计算的（顺序：compile → MSA → ICI → SparseCore → runtime → scheduler → megascale）。匹配两个组的模板只计一次，归入较早的组。总数（~2,937）带有 ±~50 的 caveat，因为 error-prose gate 会纳入少量 help strings 或排除少量短错误；这些正确归属为提示字符串（hint-strings.md），而不是错误。

交叉引用

Runtime 与执行概览 — 这些 Status 对象在 PJRT 执行路径中产生和传播的位置
提示字符串 — 建议表面（try/consider/recommended），与此处失败形成对照；±50 的边界案例位于那里
内部 Pass 名称 — compile/HLO 错误模板按名称引用的 pipeline-stage 标识符字符串
Execute Async on Stream — executor 入队路径，运行时暴露 runtime/driver StatusOr<T> 模板
Error / Status Codes（附录） — 这些模板被包装进的 absl::StatusCode 枚举
Memory-Space 表（附录） — Unsupported memory space / memory space colors 模板命名的 memory-space color 值

Error/Status 字符串模板 ​

摘要 ​

如何阅读本目录 ​

两种格式化习惯用法 ​

占位符语法 ​

置信度与归属 ​

Compile / HLO / Codegen / Verifier — ~875 个模板 ​

Scheduler / Fuel / FIFO — ~15 个模板 ​

MSA / Memory / Allocation — ~79 个模板 ​

ICI / Collective — ~90 个模板 ​

Megascale（DCN Runtime / Aggregator）— ~21 个模板 ​

SparseCore / Embedding — ~68 个模板 ​

Runtime / Driver / PJRT — ~177 个模板 ​

PJRT-C-API / protobuf-descriptor（链接库，非 TPU） ​

CHECK / RET_CHECK / 自检模板 ​

Fatal / Abort 表面 {#fatal--abort-surface} ​

Status-Code 映射 {#status-code-mapping} ​

参数类型解码（mangled type pack） ​

StatusCode 关键词分布 ​

面向用户与内部的划分 ​

概览：各子系统模板数 ​

交叉引用 ​

Error/Status 字符串模板

摘要

如何阅读本目录

两种格式化习惯用法

占位符语法

置信度与归属

Compile / HLO / Codegen / Verifier — ~875 个模板

Scheduler / Fuel / FIFO — ~15 个模板

MSA / Memory / Allocation — ~79 个模板

ICI / Collective — ~90 个模板

Megascale（DCN Runtime / Aggregator）— ~21 个模板

SparseCore / Embedding — ~68 个模板

Runtime / Driver / PJRT — ~177 个模板

PJRT-C-API / protobuf-descriptor（链接库，非 TPU）

CHECK / RET_CHECK / 自检模板

Fatal / Abort 表面 {#fatal--abort-surface}

Status-Code 映射 {#status-code-mapping}

参数类型解码（mangled type pack）

StatusCode 关键词分布

面向用户与内部的划分

概览：各子系统模板数

交叉引用