Error/Status String Templates

All offsets, symbols, and counts on this page apply to libtpu.so from the libtpu-0.0.40-cp314 wheel (libtpu_lts_20260413_b_RC00, ELF x86-64, 781,691,048 bytes, BuildID md5 89edbbe81c5b328a958fe628a9f2207d, not stripped). Other builds will differ.

Abstract

libtpu builds its diagnostics out of format-string templates baked into .rodata and filled at the call site. There are two formatting idioms and one wrapping convention. The dominant idiom is C-printf-style — absl::StrFormat("…%s…%d…", args) and LOG(...) << absl::StrFormat(...) — which accounts for the large majority of the surface; the minority idiom is absl positional substitution — absl::Substitute("…$0…$1…", args) — confined mostly to the megascale DCN runtime, the collective buffer-size validators, and a statically-linked protobuf descriptor validator. Most of these templates are then wrapped in an absl::Status by one of three factory families (see Status-Code Mapping), so the same string is simultaneously the human-readable error and the payload of a StatusOr<T> that propagates back through PJRT to a JAX/XLA user.

This page is a reference catalog, not an algorithm trace. Its value is the grouped, deduplicated table of real templates with their placeholders explained and their absl::Status code attributed where that attribution is byte-confirmed. The recovered surface is roughly 2,937 distinct error/status templates; this page does not reproduce all of them — it documents the structure of the space (the placeholder grammar, the subsystem partition, the three Status-construction idioms, the byte-confirmed argument types) and gives a representative, source-anchored sample per subsystem so a reimplementer can predict the shape of any template not listed.

Placeholder semantics are the through-line. A %s is almost never a raw C string: it is some object's ToString() / AbslStringify — a Shape, a Layout, an HloInstruction name, a device name, an opcode mnemonic. A %d/%lld/%zu is a dimension, an ordinal, a count, or a byte size. Pointers (%p) and floats (%f) appear only in low-level driver and cost-model paths. The catalog calls out, per template, what each placeholder means, because the printf type alone (%s = char*) hides the real C++ argument.

NOTE — the error/status templates here are distinct from user-facing hint strings (suggestions phrased as "try…", "consider…", not failures), which live on hint-strings.md, and from the internal pass-name strings (pipeline-stage identifiers, not diagnostics) on internal-pass-names.md. This page owns the error/status format-string catalog, the placeholder semantics, and the status-code mapping.

How To Read This Catalog

The two formatting idioms

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)

The printf surface uses %-specs; the positional surface uses $0, $1, … (and a handful of %v, the absl stringify-any sigil). The two never mix inside one template. A reimplementer keying off %-specs alone will miss the ~138 positional templates entirely.

Placeholder grammar

The full %-spec occurrence distribution across all templates is heavily skewed toward strings and ints:

GOTCHA — %s does not tell you the argument is a string in the program. It tells you the call site passed a const char*, which is overwhelmingly the product of an object's ToString(). When reimplementing, the real type behind a %s is recoverable only at the call site (e.g. whether the operand of "%s vs. %s" is a Shape::ToString() or an HloInstruction::name()); the catalog notes the likely type but does not byte-confirm it for the printf family.

Confidence and attribution

Templates marked CERTAIN were spot-confirmed verbatim in the decompiled call site. The subsystem grouping is a keyword classification of the prose, not a per-template call-site trace; it is HIGH confidence for the prose content and MEDIUM for the exact owning pass. The status code column is CERTAIN only for the 99 <Code>StrCat factories (the mangled symbol names the code) and for the 390 MakeErrorStream sites (the macro names the code); for every other template the code is inferred from the prose and the dominant idiom of its subsystem, and is marked MEDIUM or LOW accordingly.

Compile / HLO / Codegen / Verifier — ~875 templates

The single largest block. These are the shape-inference and HLO-verifier diagnostics a JAX/XLA user hits when a program is malformed: rank checks, dimension-range checks, shape-equality checks, layout mismatches. Almost all are %s/%d-heavy and almost all wrap InvalidArgument (a few RET_CHECK → Internal). The verifier's signature is the vs. comparison style — "%d vs. %d", "%s vs. %s" — which appears wherever two quantities must match.

NOTE — 0x858400c ("Bad scalar opcode in slot 0…") resolves in the decompile to platforms_deepsea::jellyfish::isa::DecoderBcsDf::DecodeScalar0Slot (and a per-gen DecoderJf twin). Its %v is the TensorCoreBundle's AbslStringify, and it is one of the 6 TensorCoreBundle-bearing <Code>StrCat factories (see Arg-Type Decode). The per-generation decoder variants (gxc/gfc, gxc/glc, vxc) each emit their own copy, so near-identical templates with the same prose are separate rodata strings.

Scheduler / Fuel / FIFO — ~15 templates

Latency-hiding scheduler, annotation-range checks, the --xla_fuel budget, and the on-device FIFO push/pop ordering. The annotation templates lean on %c for the literal brace/bracket characters; the FIFO templates chain %s :: %s to attach instruction context.

NOTE — 0xa03ca6a ("Illegal value for --xla_fuel…") was confirmed inside the xla::MakeDebugOptionsFlags flag-parsing closure — it is a flag-value validator, the error counterpart to the --xla_fuel flag documented on the flag-name side. The two "out of fuel" strings at 0x8628660/0x862868f are diagnostic LOG output, not Status payloads (no factory wraps them); they are catalogued for completeness but should not be assumed to be recoverable as a Status code.

MSA / Memory / Allocation — ~79 templates

Memory-space assignment, prefetch/alternate-memory, HBM defragmentation, and the heap allocator. Byte-size placeholders dominate (%lld, %zu). The two over-budget templates wrap ResourceExhausted; mismatch/verification templates wrap Internal.

NOTE — do not attribute the near-identical template 0xa13e8e5 ("Failed to allocate node (%zu bytes). Memory limit: %zu [bytes]. Used: %zu [bytes].)") to TPU memory-space assignment. It lives in perfetto::protovm::RwProtoCursor::CreateNodeFromField — the Perfetto tracing library's arena allocator, not the TPU MSA path. The genuine TPU heap-allocator over-budget string is 0xa1300d0 ("Failed to allocate %zu bytes…", row above), confirmed in xla::AlignedAllocator::Allocate.

ICI / Collective — ~90 templates

Inter-chip-interconnect link health, routing, GTC synchronization, and the collective (all-reduce / all-gather / reduce-scatter / all-to-all) buffer-size validators. This block mixes printf-style ICI driver errors with the positional $N collective validators (which are megascale-tagged — see the next section). Codes lean Internal and DeadlineExceeded.

NOTE — 0x8583d20 ("ICI Probe failed…") was confirmed in asic_sw::driver::deepsea::ici::SliceConfiguration::GetLocalTopology; 0x9d1493e ("ALL_TO_ALL not supported…") in xla::megascale::runtime::HostCommandSchedulerFactory…GenerateCommunicationIrsFromTransferRegistry. The collective buffer-size validators are positional ($N) even though they sit on the ICI/collective path — they are emitted from the megascale runtime, which is the positional-idiom stronghold.

Megascale (DCN Runtime / Aggregator) — ~21 templates

Cross-host data-center-network coordination: barrier-participant accounting, the corrupted-buffer detector, launch-id timeouts, and the coordinator's error digest. This is the heart of the positional $N idiom. The coordinator's hang-digest emits one prose variant per cause branch.

The coordinator's hang digest emits the prose "Megascale detects a hang that is likely caused by …" once per cause branch (BAD_TPU_CHIP, BAD_SC_CHIP, DATA_INPUT_STALL, DIFFERENT_MODULE, FINGERPRINT_MISMATCH, NETWORKING_ISSUE, PROGRAM_NOT_QUEUED, UNKNOWN_CAUSE), and the operator-actionable follow-ups ("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"). The abort path is in Fatal / Abort Surface.

SparseCore / Embedding — ~68 templates

SparseCore (xla_sc_) and BarnaCore embedding configuration: alignment requirements, table/feature counts, the partitioner objective enum, and the SMEM row-pointer budget. User-facing (a model-config error) and almost all InvalidArgument.

NOTE — 0xa0d36d1 ("Invalid num_features…") was confirmed in tensorflow::PopulateMissingFieldsInTPUEmbeddingConfig. The %f floats in 0xa11773a are a rare case where %f is genuinely a configuration ratio, not a cost-model internal.

Runtime / Driver / PJRT — ~177 templates

The driver state machine, device/ordinal validation, firmware-queue transitions, DMA-buffer accounting, and the PJRT C-API boundary. The most idiom-mixed block: %p and errno (%d) appear here, and codes split between FailedPrecondition (state-machine guards) and Internal.

NOTE — 0x8679159 ("device ordinal value (%d) must be non-negative") was confirmed in stream_executor::StreamExecutorAddressAllocator::GetStreamExecutor. Many runtime templates feed the executor's async stream path; see execute-async-on-stream.md for where these surface during enqueue.

PJRT-C-API / protobuf-descriptor (linked library, not TPU)

The positional $N family is also populated by a statically-linked protobuf extension-declaration validator. These are catalogued for completeness but are not TPU code and should not be attributed to libtpu's own surface:

CHECK / RET_CHECK / Self-Check Templates

The absl CHECK/QCHECK/DCHECK family and XLA's TPU_RET_CHECK do not carry full message templates — they emit a fixed prefix and then append the stringified source expression and any streamed << "msg". The comparison macros (CHECK_EQ/NE/GE) append the operand values with the "%d vs. %d" / "%s vs. %s" format that recurs across the verifier block.

GOTCHA — the CHECK(expr) << "msg" macro inlines the source expression text (spmem_buffer_type != nullptr, dynamic_size, nullptr) into .rodata. Large fragments of literal C++ source — even whole lambda bodies — appear in the string table as a side effect. These are CHECK-condition evidence, not error templates in the printf sense, and are excluded from the template count. A reimplementer scanning strings for "templates" must filter them out or be flooded with source fragments.

0xa183292 was confirmed in tpu::internal::RetCheckFailSlowPath; 0xa2300c5 in xla::status_macros::MakeErrorStream::Impl::~Impl (the destructor self-check that fires when a built-but-unconsumed Status is dropped).

Fatal / Abort Surface

The intentional-abort surface is small and centralized — ICI hard failures, the megascale coordinator abort, internal-bug LOG(FATAL)s, and one library guard that is not TPU-specific.

The megascale coordinator's 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." — is a LOG(FATAL) gated on a flag.

NOTE — the 0xbe7d460 family (12 variants: aes, avx, mmx, pclmul, popcnt, sse, sse2, sse3, sse4.1, sse4.2, ssse3, …) is the absl CPU-feature startup guard, a library abort that fires before any TPU code runs if the host CPU lacks an ISA the binary was built for. It is not a TPU diagnostic. 0x8a2941d ("Fatal error in creation of RWB Fusion…") was confirmed in xla::jellyfish::TpuInstructionFusion::RwbFusionHelper — a genuine internal-bug fatal with a "file a bug" marker.

Status-Code Mapping

A template becomes an absl::Status through one of three factory idioms. Only the first two name the code in the binary; the prose family and the bare printf-into-LOG paths require call-site disassembly to confirm the code.

The <Code>StrCat factory is the unique place where format-arg C++ types are recoverable without disassembly: each instantiation emits an Itanium-mangled symbol _ZN3xla<len><Code>StrCatIJ<typepack>EEC2E… whose <typepack> names every interleaved literal segment and typed argument byte-for-byte.

Arg-Type Decode (mangled type pack)

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

Worked decode (each a real instantiation symbol):

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)

Byte-confirmed arg-type frequency across the 99 factories: std::string 29, std::string_view 15, long 13, TensorCoreBundle 6, size_t 5, StatusOr<…> 3, const int& 2, float 2, TpuVersion 1. Even at the byte-confirmed level, args are overwhelmingly string-ish + long; float and pointer are vanishingly rare — consistent with the printf-spec distribution above.

StatusCode keyword distribution

String-table keyword occurrence (noise-filtered), an upper-bound prose signal of which codes the surface favors — not a per-template count:

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

GOTCHA — the raw Internal keyword count (~16,823) is dominated by the /internal/ source-path component in inlined __FILE__ strings, not by error templates. Do not read it as an Internal-status count. The byte-confirmed Internal total is the 29 InternalStrCat factories plus the MakeErrorStream Internal sites.

User-Facing vs Internal Split

The prose itself signposts the audience. User-facing (operator / JAX-user actionable, no bug id): the HLO verifier shape/rank/dimension block, the --xla_fuel flag-value error, the embedding-config errors with their valid-value hints, the megascale digest follow-ups, and the device/executable mismatch. Internal (XLA-bug markers): strings carrying b/<id>, go/<link>, "please file a bug", or "should not happen".

NOTE — b/<id> and go/<link> tokens are the strongest internal signal. They also tag known-gap TODOs ("TODO(b/157237781) support VFIO device %s", "TODO: b/475913712 - Expected Gather indices to be bitpacked …"), which are not live errors but become Unimplemented messages when their guarded path is hit.

At-a-Glance: Templates per Subsystem

GOTCHA — the subsystem counts are first-match-wins over a prose keyword set (ordering: compile → MSA → ICI → SparseCore → runtime → scheduler → megascale). A template that matches two groups is counted once, under the earlier group. The total (~2,937) carries a ±~50 caveat where the error-prose gate admits a few help strings or rejects a few terse errors — those are properly hint strings (hint-strings.md), not errors.

Cross-References

• Runtime & Execution Overview — where in the PJRT execution path these Status objects originate and propagate
• Hint Strings — the suggestion surface (try/consider/recommended), contrasted with the failures here; the ±50 boundary cases live there
• Internal Pass Names — pipeline-stage identifier strings that the compile/HLO error templates reference by name
• Execute Async on Stream — the executor enqueue path that surfaces the runtime/driver StatusOr<T> templates at run time
• Error / Status Codes (appendix) — the absl::StatusCode enumeration these templates are wrapped in
• Memory-Space Table (appendix) — the memory-space color values named by the Unsupported memory space / memory space colors templates