Symbol & Offset Index

Every address on this page is a virtual address (VMA) for neuronx_cc 2.24.5133.0+58f8de22, recovered from the cp310 wheel. cp311/cp312 carry byte-near-identical twins of every binary, but renumber addresses — treat every value as version-pinned. The per-binary VA frame is not uniform: in the .so libraries (libBIR, libwalrus, …) .text/.rodata are loaded at VA == file-offset, so nm/objdump VAs index the file directly; in the standalone hlo-opt/hlo2penguin ELFs the section base is shifted (.text file-off = VA − 0x201000, .rodata file-off = VA − 0x200000). See §0.5. Backing synthesis: D-R09 (the cross-binary type-identity matrix) plus the per-binary nm -DC dynamic-symbol tables and IDA rtti.json typeinfo sets.

Abstract

This appendix is the reverse lookup for the whole wiki: given a symbol name or a hex address cited anywhere in the book, it names the binary that owns the symbol and the page that documents it. The normative pages anchor every claim to a sub_ADDR / 0xNNNN; this index lets a reader go the other way — from an address found in nm, objdump, or a crash backtrace to the page that explains what that code does. It is curated, not exhaustive: it carries the few hundred symbols a reader will actually look up (the dispatch cores, the per-engine encoders, the allocator drivers, the typeinfo anchors), not every one of the 59,466 functions in libwalrus or the 44,794 typeinfos in hlo2penguin.

The index is organized by binary, because which artifact a symbol lives in is the first disambiguator — the backend is split across eight .so plus two standalone ELFs and a galaxy of Cython modules, and the same short name (run, visit*, schedule) recurs in several of them. §0.1 gives the binary roster and its ownership boundaries (from D-R09); §0.5 gives the two VA-framing conventions and the two-VA-frame sidecar artifact that must be understood before any address on this page is unambiguous. §1–§8 are the per-binary symbol tables. §9 reproduces the cross-binary class/namespace matrix from D-R09 — who defines each shared type and how it crosses binary boundaries.

This is a reference catalog; its value is being a faithful, navigable index. Every row points at a page that actually exists in SUMMARY.md. Where a symbol's primary page has not been written yet, the row says so rather than inventing a target. The five highest-traffic symbols were adversarially re-verified against the binary symbol tables (§10); each row carries a confidence tag.

0.1. The binary roster

The compiler's logic is split across the artifacts below. Ownership is from D-R09 — the canonical definer of each shared C++ namespace (OWN), not merely a binary that references it. Sizes and paths are catalogued on 0.4 Binary Inventory & the .so Map; this table is the per-binary index header.

NOTE — libBIR.so is the leaf / owner: its DT_NEEDED lists only system/util libraries, and it exports bir/pelican typeinfos as weak (V) symbols. Every other backend binary reaches bir/pelican by dynamic link to libBIR (DT_NEEDED + weak-symbol coalescing), not by carrying a static copy — so a bir::* typeinfo address in those binaries' rtti.json resolves to libBIR's definition at load (D-R09 §4d-i). When this index lists a bir/pelican address it is libBIR's, the definitional one.

0.5. The two VA-framing conventions

Two distinct conventions decide whether an nm/objdump VA can be used as a file offset, and a third artifact (the two-VA-frame sidecar) can make one function appear at two addresses. Resolve all three before reading any address below as authoritative.

(A) the .so libraries  (libBIR, libwalrus, libBIRSimulator, …)
      .text VA == file offset      .rodata VA == file offset
      ⇒ an nm/objdump VA indexes the file directly; jump tables read by VA.
      e.g. setupHeader @ 0x143f440 IS at file offset 0x143f440 in libwalrus.

(B) the standalone ELFs  (hlo-opt, hlo2penguin)
      .text  file_off = VA − 0x201000
      .rodata file_off = VA − 0x200000
      ⇒ a .text VA must have 0x201000 subtracted before xxd/objdump --start-offset.
      e.g. NextChannelId @ 0x8ab1ac0 (VA) is at file_off 0x8aaa0c0 (.text) in hlo-opt.

(C) the two-VA-frame sidecar artifact  (any binary with per-symbol decompile)
      The ELF symbol-table VA (cited HERE) is the authoritative frame.
      The per-symbol IDA/Hex-Rays sidecar may carry a SECOND internal frame
      (a thunk wrapper) — e.g. is_valid_s3_lw @ 0x14ac170 (symtab) vs 0x5fbd40
      (sidecar thunk). BOTH denote the same function; a sidecar's 0x5fbd40 is
      NOT a wrong address. The .plt thunk band (libwalrus 0x5e9020–0x62d650) is
      this same effect: each thunk tail-jumps to a real body > 0xf00000.

GOTCHA — the convention is per binary, not global. Over-generalising "VA == file offset" onto an hlo-opt/hlo2penguin address is off by 0x201000/0x200000, and onto a .data-resident struct is off by the .data VMA/file-offset delta (a separate 0x400000-class shift not used by any .text/.rodata symbol below). Every address in this index is the symbol-table VA in frame (A)/(B); when a page cites a smaller sidecar number for the same symbol, that is frame (C), not a contradiction.

1. libBIR.so — the BIR core

The definitional root: bir::Instruction and the 110-opcode hierarchy, the pelican symbolic algebra, the Dtype tables, the cast engine, and the BIR-JSON SerDe. .text/.rodata VA == file offset.

Core IR & opcode enum

pelican symbolic algebra

Cast engine & BIR-JSON

2. libwalrus.so — the backend

The 65 MB backend: ~150 passes, the dependence graph, the three schedulers, the SB/PSUM/DRAM/Reg allocators, loop optimization, the LNC splitter, lower_dma, the CoreV{2,3,4}Gen per-engine code generators, the BIR linker, NeffPackager, the DebugInfoWriter, the legality cascade, and getArchModel. .text (0xf00000+) / .rodata (0x1c72000+) VA == file offset. Retains its full dynamic symbol table (nm -DC is authoritative).

Arch model & pass registry

Per-engine code generators (the L3 wire path)

Schedulers & allocators

Legality cascade

KLR→BIR codegen (beta2) & NEFF packaging

DebugInfoWriter (the provenance producer)

NOTE — DebugInfoWriter is the sole producer of the neuroncc::debug_info protobuf (D-R09 §2c). Its six core classes are byte-identical in libBIR, libwalrus, and hlo2penguin, but they cross the frontend↔backend cut by independent compilation of the same generated .proto, not by a link — hlo2penguin links none of the backend .so. This is the single cross-cut domain type (§9 row neuroncc::debug_info).

3. libBIRSimulator.so — the functional simulator

The BIR functional simulator (birsim), the whole-machine state model, and the DAP debugger half. Owns the engine half of the birsim namespace (10 disjoint classes: Memory, SyncState, TpbEngines, NeuronCoresManager, PhysicalMemory, BaseLogger) — disjoint (shared=0) from libwalrus's birsim::debugger half (D-R09 §2e/§5-C2). Links libBIR (dynamic) + libpwp_sim. .text/.rodata VA == file offset.

QUIRK — birsim is a split namespace. The debugger/driver half (39 classes, DebuggerCallbacks @ 0x1109430) lives in libwalrus; the engine half (10 classes) lives here. A reimplementer looking for birsim::Memory in libwalrus will not find it, and vice-versa for birsim::debugger. The two share zero typeinfos (D-R09 §5-C2 corrects the earlier "libwalrus-only" framing).

4. hlo-opt — the HLO optimizer

The standalone HLO/MHLO/StableHLO optimizer ELF and its --passes registry. Frame (B): .text file_off = VA − 0x201000, .rodata file_off = VA − 0x200000. Statically links upstream OpenXLA/MLIR/LLVM (D-R09 §3) — the Neuron-authored passes are a thin layer over that bulk.

Collective stream/channel-id family

CORRECTION (audit #820) — the masks 0x650 / 0x810000000000001 belong to NextChannelId @ 0x8ab1ac0 and the inlined checker #58, not to IsCollective @ 0x1f7e010. IsCollective uses 0x20000410 / 0x80011 (high band indexed by op − 0x53) plus op == 7, selecting opcodes {0x4, 0x7, 0xa, 0x1d, 0x53, 0x57, 0x66}. The three predicates are separate; their masks must not be cross-applied. (Carried verbatim from 4.4 collective-stream-channel-id.)

Pass registry

5. hlo2penguin — the MLIR front-half

The standalone MLIR front-half ELF: mhlo/stablehlo → Penguin Python IR. Frame (B) (same −0x201000 / −0x200000 shift). 44,794 unique typeinfos, of which ~0.5 % are Neuron-authored — the rest is statically-linked upstream MLIR/XLA/LLVM/oneDNN, with oneDNN present but dead on the Trainium path (D-R09 §2g/§4a/§5-C6). Links none of the backend .so (D-R09 §3-B1).

NOTE — mlir (10,328), xla (4,634), dnnl (13,786 substring / 2,100 _ZTI), absl (1,899), and llvm (7,033) typeinfos are upstream, statically linked — not a Neuron domain boundary. They are deliberately not indexed per-symbol here: a reimplementer wanting mlir::* reads upstream MLIR, not this wiki. The oneDNN bulk is dead code on the Trainium path (D-R09 §5-C6).

6. libpwp_sim.so — the PWP simulator

A fully isolated domain (D-R09 §2h/§5-C5): PWPSim::Simulator + PWPSim::AFTable, 4 types total. NEEDED by libwalrus/libBIRSimulator/nki_klr_sim for the numeric activation model, but exposes no shared bir/pelican/klr type. .text/.rodata VA == file offset.

7. libBIRParserDumper.so — the BIR-JSON round-trip

parserdumper / bir_roundtrip — the BIR-JSON read-write round-trip serializers. Links libBIR (dynamic). 557 typeinfos. Carries 26/29 of the pelican classes and a bir subset + its own bir::IRVisitor<DumperPass> instantiations (D-R09 §1).

8. The Cython codegen modules

The NKI forward builder, the beta3 Penguin→BIR driver, and the re-emit printer ship as per-Python-version Cython extension .so. Frame (A) (VA == file offset for .text/.rodata); the __pyx_pw_* / __pyx_pf_* symbols are the Cython-mangled method labels. KernelBuilder.so ships unstripped with DWARF — the most readable binary in the NKI stack.

QUIRK — KernelBuilder.so/NeuronCodegen (the forward builder, NKI Python → Penguin IR) and NkiCodegen.so/NkiCodegen (the reverse printer, Penguin IR → NKI text) are different classes in different files. They are not the two ends of one pass — NkiCodegen is the inverse direction. Confusing them is the most common NKI-layer error; see 0.4 binary-inventory.

9. Cross-binary class/symbol matrix (D-R09)

The ownership matrix from D-R09 §1 — namespace (row) × binary (col). OWN = canonical definer (weak V typeinfo); dyn = present by dynamic link to the owner (DT_NEEDED + weak coalescing); LOCAL = binary-local, not shared; · = absent (CONFIRMED zero hits). Numbers in () are distinct-type counts measured in that binary's rtti.json. This is what disambiguates which binary's copy a typeinfo address refers to.

* — neuroncc::debug_info in hlo2penguin is not reached by a link (it links no backend .so); the 6 classes are the same generated protobuf compiled independently into each binary. The 8 _ZTI names are byte-identical across libBIR ∩ libwalrus ∩ hlo2penguin — the only frontend↔backend crossing type.

The domain boundaries (D-R09 §3): frontend ↔ backend is a clean cut — hlo2penguin has bir=0, pelican=0, klr=0, dap=0, birsim=0 (absolute zero), and the backend .so have xla=0, mlir=0, dnnl=0. The single exception is the six neuroncc::debug_info classes. The DT_NEEDED DAG is rooted at libBIR (leaf) → libBIRParserDumper/libBIRSimulator → libwalrus (top), with libpwp_sim isolated and hlo2penguin linking only {libm, libstdc++, libgcc_s, libc}.

10. Adversarial verification — five highest-traffic symbols

Each was re-read from the binary's own dynamic symbol table (nm -DC) and the cited page re-checked for existence. All five resolved exactly.

NOTE — visitInstCustomOp is the canonical example of why this index is per binary: the name appears at 0xfa7bd0 (birverifier::InstVisitor, the verifier — 8.41), at 0x12613c0 (CoreV2GenImpl, the encoder — 11.8), and as an undefined U (birsim::InstVisitor, defined in libBIRSimulator — 7.39). A bare "visitInstCustomOp @ …" is ambiguous until the binary and the address fix which one.

Cross-References

• 0.4 Binary Inventory & the .so Map — the artifact atlas this index keys on: sizes, paths, the five distinct ~230 MB tool ELFs, the eight .so, the Cython galaxy.
• 0.6 Build & Version Provenance — the cp310/311/312 cross-wheel parity argument and why addresses drift across wheels.
• 14.1 Master Opcode Reference — the opcode-space sibling appendix; setupHeader / InstructionType2string / enum_variant_string_opcode are documented there in full.
• 14.2 Master Dtype / Enum Reference — the dtype/enum-ordinal sibling appendix.
• 1.1 The Arch Object Model — getArchModel and the Board/Device/Core accessor chain.
• 7.27 KlirToBirCodegen Dispatch Core — the 77-case master routing table behind codegenStmt.
• 8.39 The Opcode Master — the setupHeader L3 wire path.