Confidence Ledger
This page indexes claims and gaps across the
neuronx_cc2.24.5133.0+58f8de22 wiki (cp310/cp311/cp312 wheels, static analysis only). Every entry below quotes the confidence tag as it appears on the source page; the addresses and symbols belong to the cited page and are pinned there. When a source page is re-graded, this ledger must follow — it is a mirror, not an independent claim.
Abstract
Every page in this book carries the four-tier confidence ladder from §0.3: CONFIRMED (directly observed), STRONG (multiple signals agree, decisive line not read end-to-end), INFERRED (deduced from structure or a confirmed sibling), SPECULATIVE (a plausible unproven reading). Most of the book is CONFIRMED or STRONG. This appendix is the opposite-facing index: it consolidates, in one place, the book's weakest-evidence claims and its genuine gaps, so a reader can see exactly what is not solidly nailed down without re-reading 200 pages to find the caveats.
The honesty argument is the whole point. A wiki that hides its soft spots is more dangerous than one that has none, because a reader cannot tell a confident claim from a hopeful one. §0.3 already lists the categories of thing that are provably not recoverable from the shipped binaries (Xtensa instruction bodies, the SUNDA_APB_BASE numeric, identifier strings that were never compiled in). This page is the instance-level counterpart: the specific pages, the specific claims, the specific reason each is soft, and — for each — the evidence that would settle it.
Entries are grouped by why they are weak, not by topic, because the reason is what a reader needs in order to weigh the claim:
- Structure-only — no code bodies were recoverable; the claim rests on ELF headers, byte-diffs, and strings (the GPSIMD Xtensa images, §11.1).
- Cython-obscured call-order — the symbol or string is confirmed, but the body lives in a compiled
.sowhose control flow was not walked instruction-by-instruction (much of Part 6,parallel_state,CollectiveOp.so). - DWARF-reconstructed — the body was read, but via the
KernelBuilderdebug line table rather than disassembly; the line→method map is the evidence, not the opcodes. - Cross-library / PLT-thunk — the call site is pinned but the callee body is in a different binary (
libBIR.so,parallel_state) not walked here. - Genuine gaps — a value, struct, or path that is simply not in the corpus: a withheld source, a stripped packer, an Xtensa immediate, a config-supplied default.
- Flagged seams — a claim graded STRONG, not CONFIRMED on its own page because the decisive plumbing crosses a boundary (the front↔back LNC-size equality; the
info.jsonctor seed).
The §-numbers below are this book's part/chapter scheme; the link is the canonical page. Where a page grades a claim, the exact tag string is reproduced so this ledger can be audited against it (see §7).
| Scope | The wiki's own confidence tags, consolidated |
| Confidence model | §0.3 Methodology & Confidence |
| Strongest single gap | GPSIMD Xtensa code — no disassembler in corpus (§11.1) |
| Most-cited soft seam | front-end logical_nc_config == backend lnc_size — STRONG, not CONFIRMED |
| Tag vocabulary | CONFIRMED / STRONG / INFERRED / SPECULATIVE / GAP |
| Maintenance rule | mirror only — re-grade the source page first, then this row |
1. Structure-only — claims with no recovered code body
These pages document hardware or images whose instruction bodies could not be disassembled at all. Every behavioural claim is INFERRED from ELF structure, a cross-image byte-diff, and .rodata/.data strings; the layout claims are CONFIRMED. The pages say so plainly and never present a string-derived behaviour as observed code.
| Claim | Page | Tag on the page | Why structure-only | What would resolve it |
|---|---|---|---|---|
GPSIMD CPUs are 8 Tensilica Xtensa ELF32 images, one per core, re-linked at 0x84000000 + id·0x200000 | §11.1 | CONFIRMED (layout) | file/readelf + byte-diff prove it without code | already settled |
cpu_id derivation reads MEM_WINDOW0_LO UREG and checks SUNDA_APB_BASE | §11.1 | STRONG (string + byte-diff) / SUNDA_APB_BASE numeric SPECULATIVE | the deriving code is Xtensa, not disassembled; the immediate is an Xtensa literal | an Xtensa disassembler, or a NEFF/trace fixture showing the resolved base |
The 6 .ctors register ATen/c10 statics + op-name table | §11.1 | count CONFIRMED; binding STRONG (from .rodata strings) | ctor bodies are Xtensa, not read | Xtensa disasm of the ctor entries |
| Op dispatch is a flat function table keyed by FunctionId, not a c10 Dispatcher | §11.1 | STRONG (string surface; Dispatcher symbols absent) | absence-of-symbol argument, not a read selector | Xtensa disasm of the entry trampoline |
| Headroom (~1.37 MiB/core) holds stack/heap/DMA scratch | §11.1 | span CONFIRMED (arithmetic); use INFERRED | no allocator code disassembled | Xtensa disasm of the runtime allocator |
HARD LIMIT (from §11.1) — host binutils has no Xtensa backend (
objdump -i | rg xtensa→ empty) and IDA recoveredtotal_functions=2, decompiled=0, flirt=null. Any future page claiming an Xtensa instruction sequence must ship the disassembler that produced it. The companion reconciliation page GPSIMD reconciliation is not in this row — its compiler/IR side (libwalrus/libBIR) is disassembled; only its cp312 anchor0x13597a0is ASSERTED because cp312libwalrusis not in the indexed corpus.
2. Cython-obscured — symbol confirmed, body not walked
The NKI front-end (Part 6) and the distribution layer (Part 13) lean heavily on compiled Cython .so modules. Their string surface (class rosters, method names, pyx_n_s_ identifiers, asserts) is binary evidence and is CONFIRMED. What is soft is the control flow inside those modules — the exact call order, the default geometry, the field offsets — when the body was not single-stepped.
| Claim | Page | Tag on the page | Why soft | What would resolve it |
|---|---|---|---|---|
| Default worker-group geometry behind the 5 group collectives | §6.5.x NeuronCodegen Collectives | INFERRED (group geometry) | gen_all_worker_group body lives in parallel_state, a module not walked here | disasm of parallel_state.gen_all_worker_group |
CollectiveKind numeric resolution from Op class | §6.5.x / §6.5.13 | roster CONFIRMED; the resolve step STRONG | CollectiveOp.so name pool is confirmed; the Op→kind mapping body is downstream | disasm of CollectiveOp.so resolver |
The Shardy export (dims, reshape_dims, transpose_perm) perm arithmetic | §13.x Shardy↔HloSharding bridge | INFERRED — "call order CONFIRMED, perm vector not traced bb-by-bb" | the 279-bb convertToHloSharding was not traced block-by-block | bb-by-bb trace of convertToHloSharding @0x2bc58f0 |
NOTE — the one Cython module that is not obscured is
KernelBuilder.cpython-3xx.so, which ships with full DWARF debug info (§0.3 QUIRK). Pages that lean on it are in §3 below, not here, because there the body was read — through the line table.
3. DWARF-reconstructed — body read via the debug line table
KernelBuilder.cpython-3xx.so is the most readable binary in the stack: its DWARF line table maps every method to KernelBuilder.py line ranges. The NKI codegen pages (Part 6) use this to recover method sequence and call sites. This is strong evidence — stronger than string-only — but it is worth flagging that the underlying anchor is a <file>:<line> map, not a disassembled opcode stream: a :NNN source line cannot be re-verified against a stripped peer .so, and the line table reflects the source structure, not necessarily the emitted control flow after optimization.
| Claim class | Pages | Tag posture | The seam |
|---|---|---|---|
KernelBuilder method call orders (KB.py:NNNN anchors throughout) | §6.5.x NeuronCodegen* pages | CONFIRMED at method granularity via DWARF | the :NNNN line is from the debug table; trust the method, spot-check the order |
info.json pre-seed in NeffFileWriter ctor | §12.x NEFF header writer | downgraded to STRONG / INFERRED; source-line anchor 0x1543eb0:228 dropped | a :NNN line cannot be verified against a stripped .so; the seed is corroborated (the string is loaded in the ctor) but not single-stepped to the insert call |
4. Cross-library / PLT-thunk — call site pinned, callee elsewhere
A recurring soft spot: the wiki pins a call site in libwalrus.so but the callee body lives in libBIR.so (imported as a PLT thunk) and was not traced. The result use is certain; the internal decode is taken on faith from a named model.
| Claim | Page | Tag on the page | Callee location | What would resolve it |
|---|---|---|---|---|
EngineAccumulationType enum→bool decode (getCalcStart/getCalcAccu) | §7.x Sim MatMul-MX | GAP (G1) — result-use certain, bit-decode taken on faith | libBIR PLT thunks @0x22b0a80 / @0x22b1618 | disasm of the libBIR accessors |
Per-function key 19 ("auto psum accumulate") authoring pass | §7.x | GAP (G5) — sim reads it; the setter is a libwalrus/HLO concern | upstream authoring pass, not traced | trace the pass that writes key 19 |
DynamicAPINFO setter bodies | §5.x Symbolic AP register-ALU | GAP | defined in libBIR, imported here | disasm of the libBIR setters |
5. Genuine gaps — values and paths simply not in the corpus
These are not "soft readings" — they are absences. The value, struct, or source does not exist in the shipped artifacts in any readable form. A reimplementer should treat each as a hole to fill from a fixture, a header, or hardware, not from this wiki.
5.1 Source-withheld NKI production leaves
The deepest gap in the book. §6.6.4 Production Kernel Inventory documents the three-tree topology: nkilib (open .py), _pre_prod_kernels (readable glue), and _private_kernels/ — 34 compiled Cython .so, zero .py, source withheld.
| Gap | Page | Tag on the page | What would resolve it |
|---|---|---|---|
Algorithm bodies of attention, qkv, mlp, expert_mlps, router_topk, conv, collective_matmul, fused_linear, prefix_caching_attention, hw_ubench | §6.6.4 | entry-point symbols recoverable; algorithm bodies NOT ("no readable twin … short of decompiling the Cython .so") | Cython-.so decompilation, or the upstream .py |
(blockwise_mm, llama3_transformer are not in this gap — their production copy is the readable _pre_prod .py, merely Cython-compiled) | §6.6.4 | recoverable — do not list as withheld | n/a |
5.2 AllToAll device lowering
The SPMD emitter draws an explicit boundary: it emits stock GSPMD kAllToAll HLOs, and the device-side rewrite/lowering "is where this page hands off" (§13.x SPMD collective emission). The AlltoAllOp class and the AllToAll CollectiveKind name are confirmed in CollectiveOp.so, and the assert Illegal AlltoAll without {split,concat}_dimension is read (§6.5.x) — but the device lowering body (split/concat-dimension realization to BIR DMA) sits in downstream xla::hilo rewrite passes documented elsewhere, and is the genuine seam between emission and codegen.
| Gap | Page | Tag posture | What would resolve it |
|---|---|---|---|
| AllToAll → device DMA realization (split/concat dim handling) | §13.x emission boundary + §6.5.x | class/kind CONFIRMED; lowering body handed off, not traced on these pages | trace the xla::hilo AllToAll rewrite + the BIR DMA emit |
5.3 Pipeline-parallel / MPMD partition
Boundary markers for the per-layer pipeline cut are CONFIRMED as paired Start/End kCustomCall (0x2B) sentinels carrying boundaryCount=<N> (§4.x boundary markers). What is not in the corpus is an end-to-end MPMD / pipeline-parallel partition driver that consumes those cuts into separate device programs — the markers are structural metadata "stripped before codegen", and the multi-program scheduler that would act on them is not a traced pass here.
| Gap | Page | Tag posture | What would resolve it |
|---|---|---|---|
| Pipeline-parallel / MPMD multi-program partitioner | §4.x | marker mechanism CONFIRMED; the consuming MPMD driver not traced / absent | a pass reading boundaryCount to cut programs, if one ships |
5.4 Stripped-packer and config-supplied values
| Gap | Page | Tag on the page | What would resolve it |
|---|---|---|---|
generateInstLoadActFuncSet IT6 wire packing (size<<16)|((23−size)<<11)|base | §9.x BKT-ctrl blob | GAP — proven from bytes + consumer, not the packer (lib stripped) | the producing lib, un-stripped |
MaxCceDmaSource numeric (per-arch EngineInfo/Target ctor, not JSON) | §8.x DMA legalization | GAP — only the deref path is pinned | disasm of the per-arch Target ctor |
DMAQueueAttribute enumerator names (num_queues/sync_type/priority_class) | §8.x DMA queues | GAP — written structurally; names unrecoverable | a def.json attribute table or un-stripped enum |
| MX 8-partition scale-block intra-quadrant packing | § MX matmul legality | inferred from two rules; "not independently confirmed from a NEFF fixture (G4)" | a NEFF MX fixture |
sendrecv-to-gpsimd-max-bpp knob default | §8.x local collectives | GAP / SPECULATIVE (the default specifically) | the knob's per-arch default ctor |
6. Flagged seams — STRONG-not-CONFIRMED boundary crossings
These claims are well-supported but explicitly not promoted to CONFIRMED on their own page, because the decisive plumbing crosses a binary or strand boundary and was not traced through a single symbol. They are the most important rows to know about, because they read as solid and are almost solid.
| Seam | Page(s) | Exact grade on the page | The boundary |
|---|---|---|---|
front-end logical_nc_config == backend lnc_size is one quantity N | §13.8 LNC sharding constraint, §1.x LNC memory model | "STRONG, not CONFIRMED" / INFERRED-STRONG | "CLI attribute to PassOptions+0x1A4 spans the driver→penguin boundary and is not traced through a single symbol" |
| ADDR4 bit-31 register-mode flag re-purposes byte 0 as an 8-bit register id | §2.x ADDR4 | CONFIRMED (the flag + packing); name of mode bit 30 (ACTIVE) INFERRED | the register-mode flag itself is CONFIRMED (mov byte[rbx],r15b @0x1508fca); the soft part is the mode-nibble names at bits 29–30, graded "CONFIRMED (consumer); INFERRED (name)" — and the CORRECTION that bit 30 (ACTIVE) is never set by the encoder, only consumed |
info.json ctor pre-seed (also §3) | §12.x NEFF header writer | STRONG / INFERRED | corroborated by the string load in the ctor, not single-stepped to the insert call |
lnc_size front-end default (2 on Trn2/sunda, else 1) seeds the SPMD mesh | §1.x LNC memory model | STRONG | "the CLI→PassOptions plumbing is not traced through one symbol on this page" |
DependenceEdgeT member layout | §5.x backend dependence distance | INFERRED / GAPS | only the 40-byte size + the stored Instruction* are recovered |
QUIRK — a CORRECTION on a page is a confidence event, not an error. The ADDR4 page's "bit 30 (ACTIVE) is never set by the encoder" and the NEFF writer's "downgraded from CONFIRMED to STRONG" are the system working: a prior over-grade caught and demoted. This ledger lists the current grade; the page's CORRECTION block is the audit trail.
7. Adversarial self-verification
Per §0.3, this ledger is itself a claim and must not misrepresent a CONFIRMED claim as weak (or the reverse). Five rows re-checked against their source page:
- GPSIMD Xtensa "no recoverable code bodies" — §11.1 line 9 states verbatim: "there are no recoverable Xtensa code bodies … IDA's auto-analysis recovered only 2 trivial functions with 0 decompiled bodies." The §1 row is faithful: structure CONFIRMED, behaviour INFERRED. ✓
- 34
.so, source withheld — §6.6.4 line 13: "34 compiled Cython.sofiles, zero.py, source withheld." The §5.1 row reproduces the exact roster split (10 withheld,blockwise_mm/llama3_transformerrecoverable) from line 134. Not over-stated: the recoverable two are explicitly excluded. ✓ - front↔back LNC equality is STRONG-not-CONFIRMED — §13.8 line 275: "This equality is STRONG, not CONFIRMED." The §6 row quotes it verbatim and does not upgrade it. ✓
- ADDR4 register-mode flag is CONFIRMED — §2.x ADDR4 line 21 grades the bit-31 register-mode flag CONFIRMED; line 197 shows the packing store
@0x1508fca. The §6 row correctly keeps the flag CONFIRMED and flags only the mode-nibble names (bits 29–30) as INFERRED — it does not falsely demote the register-mode bit. ✓ - Shardy perm arithmetic is INFERRED — §13.x line 631: "INFERRED — call order CONFIRMED, perm vector not traced bb-by-bb." The §2 row reproduces both halves (order confirmed, vector inferred). ✓
No row in this ledger downgrades a CONFIRMED claim or upgrades a SPECULATIVE one. Where a page carries a CORRECTION, the row reflects the post-correction grade.
See also
- §0.3 Methodology & the Confidence Model — the four-tier ladder this page mirrors, and the category-level "not recoverable" catalog.
- §11.1 The GPSIMD CPUs: 8-core Xtensa ELF Layout — the canonical structure-only page (§1 here).
- §6.6.4 Production Kernel Inventory: the Three-Tree Story — the source-withheld leaves (§5.1 here).
- Report ↔ Page Cross-Reference — which backing report each page cites; this ledger audits grade, that page audits linkage.