Glossary
Version pin: every symbol/struct/address cited here is from the runtime artifacts this book derives from —
libnrt.so.2.31.24.0(aws-neuronx-runtime-lib 2.31.24.0-0b044f4ce, build-id8bb57aba0fb2e0035f1d88e9fc4fb3e7387c102e, ELF64, unstripped, DWARF v4), the GPL kernelaws-neuronx-dkms 2.27.4.0(neuron.ko), and the two firmware carrierslibncfw.so(a98f8e1c…) /libnrtucode_extisa.so(7bb03bc4…). Each term's canonical definition is owned by the page in the Owning page column; this page does not re-derive — it indexes. · Part 0 — Reference Apparatus / REFERENCE · back to index
Abstract
This is the controlled vocabulary of the Neuron runtime: one row per term, each defined from the established usage on its owning page, pinned to a canonical symbol or struct and an exact wiki/src/… path. It exists so that "vNC", "encd", "STPB", or "S2M" mean exactly one thing across the whole book — when a page writes a term, this glossary is where the single authoritative definition and the page that owns it both live. Read a row to learn what a token is; follow the owning page to learn how it works. Every definition here is derived from binary/GPL analysis; nothing is invented, and where a term's meaning is contested by two pages the disagreement is noted as a CORRECTION on the owning page, not papered over here.
The vocabulary has a structural trap a reader must internalize before using it: the Neuron stack reuses near-identical names for distinct objects at adjacent layers, and a transposition is silently wrong rather than loud. A physical NeuronCore (NC) is not the virtual/LNC-collapsed core (vNC); the sync-processor (SP/TopSP) is not the state-programming layer (STPB); the collective composer engine (enc/) is not the device-side descriptor emitter (encd/); the input ELF (kelf) is not the lowered binary (kbin) that kelf2kbin produces; memory-to-stream (M2S, Tx) is the opposite direction from stream-to-memory (S2M, Rx); the SDMA engine band is not the AnnapurnaLabs UDMA controller it programs; and the three collective algorithm families (mesh / metaring / kangaring) are not interchangeable. The alphabetical table comes first; the Confusable Pairs section that follows is the highest-value part of this page — one GOTCHA per pair, each with the distinguishing symbol and page.
The Controlled Vocabulary
Alphabetical by term. Canonical symbol/struct is the single anchor a reimplementer keys on; Owning page is the exact SUMMARY path that owns the definition; Conf is confidence in the term↔definition binding as established by the owning page.
| Term | Definition | Canonical symbol / struct | Owning page | Conf |
|---|---|---|---|---|
| basis-points (bps) | Fractional-share unit used by the GPSIMD DGE priority/QoS maps and KAI-style fractional gating — 1 bps = 1/10000; appears as the bitmask granularity in the ext-ISA core->kind and DGE priority-map programming. | core+0x10 kind bitmasks (0x02020204 / 0x20202040) | gpsimd/overview.md | MEDIUM |
| CAYMAN | Silicon codename for NeuronCore-v3 = Trainium2 (Trn2); 8 NeuronCores, arch enum value 3 across kernel/KaenaHal/encd. | AL_HAL_TPB_ARCH_TYPE_CAYMAN = 3 | arch/coretype-numbering.md | HIGH |
| CRWL | Per-NeuronCore Cooperative RW Lock in the kernel driver — a hand-rolled writer-preferring reader/writer lock (mutex + usleep_range spin, hard retry cap, no wait queue) that serializes model-load (writer) against in-flight inference (readers), keyed on a 32-byte model UUID + writer tgid. | struct neuron_crwl (nd->crwl[8]) | kernel/crwl.md | HIGH |
| DGE | Descriptor-Generation-Engine — the Pooling/Q7 engine's on-TPB hardware descriptor generator; programmed via a per-arch carveout cap (0x30000/0x38000/0x40000) and DMA-mapping pack in the STPB Pooling leaf. | pooling_set_dge_carveout (STPB leaf) | runtime/arch-stpb.md | HIGH |
| DHAL / KaenaHal | The per-generation device hardware-abstraction layer. Kernel-side DHAL is the ndhal vtable-of-vtables installed per arch (ndhal_register_funcs_v{2,3,4}); userspace KaenaHal (vendored KaenaHal-2.31.0.0, al_hal_*/aws_hal_*) is the kaena_khal HAL dispatched off the latched arch type. | kernel struct ndhal · userspace kaena_khal (.bss @0xCAEB80) | kernel/dhal-core.md · runtime/hal-adapter.md | HIGH |
| dmem | The runtime's userspace device-memory allocator — a bookkeeping/marshalling slab over the kernel MEM_ALLOC / MEM_COPY_ASYNC ioctls (via NDL); owns the 192-byte dmem_t handle (kernel mem-handle + resolved physical address) that backs every weight, tensor, ring buffer, and scratchpad. | dmem_t (192 B) · dmem_alloc_internal @0x228640 | runtime/tdrv-dmem.md | HIGH |
| enc | The host-side collectives composer engine (enc/enc.cc, the enc_*/enc_primitive* C++ classes) — turns a parsed replica group into a per-NeuronCore ring/mesh/hier schedule; statically embedded in libnrt.so (not in libnccom). | enc_context · enc_post_operation @0x11f790 | collectives/engine-core.md | HIGH |
| encd | The device-resident descriptor emitter (tdrv/encd.c, the encd_* driver) — lowers each composed enc event into the wire artifacts silicon consumes: TopSP SPAD op-streams and SDMA vring DMA packets. Called down into by enc; never calls back up. | encd_context (624 B) · encd_dma_mark_end @0x237200 | collectives/encd-overview.md | HIGH |
| FW-IO | "Firmware I/O" — the kernel↔firmware request/response mailbox in a MiscRAM aperture of BAR0; CRC32C-framed, sequence-numbered messages to the device's Q7 management CPU (not the NCFW sequencer, not the GPSIMD vector cores). | struct fw_io_ctx · fw_io_execute_request (neuron_fw_io.c) | kernel/fw-io.md | HIGH |
| genpool | The kernel HBM allocator's per-channel/region gen_pool grid (MAX_DRAM_CHANNELS 4 × MAX_DDR_REGIONS 4); device genpool VA = physical address OR'd with bit-63 (GENPOOL_DEVMEM_BASE) so pa==0 is representable. | GENPOOL_DEVMEM_BASE 0x1ull<<63 | arch/memory-hierarchy.md | HIGH |
| GPSIMD / Q7 | The on-device vector co-processor "pool" engine — a Cadence/Tensilica Vision-Q7 (Xtensa LX + IVP), programmed per-custom-op by microcode shipped in libnrtucode_extisa.so; firmware self-labels "Q7". | nrtucode_* API · core handle (0x70) | gpsimd/overview.md | HIGH |
| IVP | "Instruction Vision Processing" — the Vision-Q7's 512-bit vector ISA family (ivp_* mnemonics / IVP_* builtins); 1065 ISA-wide ops, 285 emitted by the shipped pool kernels. | Iclass_IVP_*_args (1065) | gpsimd/ivp-isa-catalog.md | HIGH |
| kangaring | A metaring_type value (=1) — Neuron's multi-rail ring variant (per-VNC kangaring_ranks[32] / logical_path[256]); a ring-family member, gated on vcore_size+rank_n. | alg_kangaring_init @0xfe480 · metaring_type=KANGARING | collectives/algorithm-taxonomy.md | HIGH |
| KBIN | The per-subgraph lowered device-binary descriptor record (424 B, in-RAM only) that kelf2kbin produces from the parsed kelf working set; one per subgraph == one per NeuronCore. There is no monolithic .kbin file on disk. | kbin (424 B, ord 7153) · gen_kbin @0x4af930 | neff/kbin-structs.md | HIGH |
| kelf | The KELF descriptor — the JSON manifest inside a NEFF that names per-subgraph def.json definitions and <eng>.bin blobs; the parse input to lowering (kelf/kelf.cpp, kelf::kelf::load). | kelf::kelf::load @0x497dc0 | neff/kelf2kbin.md | HIGH |
| kelf2kbin | The lowering transform (kelf/kelf2kbin.cpp) that turns a kelf-parsed mla_resources working set (440 B) into the in-RAM 424-byte kbin device-binary record — the JSON→KBIN compile step (NEFF load stage S3). | kelf2kbin.cpp · gen_kbin @0x4af930 | neff/kelf2kbin.md | HIGH |
| LNC | "Logical NeuronCore" — the runtime abstraction that collapses N physical NeuronCores into one virtual core (a vNC) so a model spans multiple cores as a single logical unit; MAX_VIRTUAL_TPB bounds the collapse. | virtual_core_t (physical_core_t+0x08) · MAX_VIRTUAL_TPB | runtime/tdrv-lifecycle.md | MEDIUM |
| M2M | Memory-to-memory — the UDMA copy mode/builder used for device-DRAM↔device-DRAM and host bounce-buffer copies; the leaf the vring copy-append delegates to. | al_udma_m2m_build_copy_descriptor @0x45cca0 | kernel/udma-m2m.md | HIGH |
| M2S | Memory-to-stream — the UDMA Tx descriptor class (host/device memory → stream); the al_sdma_m2s_build_* op-word + descriptor packers feed the stream side. | al_sdma_m2s_build_*_descriptor (0x451690…) | dma/meta-ctrl-overlays.md | HIGH |
| MARIANA | Silicon codename for NeuronCore-v4 = Trainium3 (Trn3, mesh topology); 8 NeuronCores, arch enum value 4. (MARIANA_PLUS = v4+, a firmware-only microcode tier with no kernel arch value.) | AL_HAL_TPB_ARCH_TYPE_MARIANA = 4 | arch/coretype-numbering.md | HIGH |
| mem_ref | The static memory-placement plan carried from compiler to silicon: every tensor gets a fixed device physical address once at load, and the executor only resolves (var, offset)→addr by an add (zero per-inference allocation). Two objects share the name — the build-side C++ tree and the load-side mem_ref_t POD. | build C++ mem_ref tree · load mem_ref_t (208 B) | neff/memory-planning.md | HIGH |
| mesh | Collective algorithm family — a fully-peered N·(N-1) exchange the ring cannot express; the only family that can serve ALLTOALL. Composed by alg_mesh_init; device pattern ENC_PATTERN_MESH=1. | alg_mesh_init @0x135990 · enc_alg_mesh | collectives/algorithm-taxonomy.md | HIGH |
| metaring | The ring-family schedule object and topology sub-axis (metaring_type: RING/KANGARING/SINGLE_CYCLE_RING/RDH) — the one composer (enc_metaring_primitive) shared by all ring-family algorithms; init_metaring_algorithm dispatches it to ring/kangaring/rdh init only. | enc_alg_metaring · init_metaring_algorithm @0xfe570 | collectives/algorithm-taxonomy.md | HIGH |
| MLA | "Multi-Logical-Accelerator" — the runtime's per-Neuron-device object (mla_t, one of 32 inline in tdrv_ctx_0), each holding eight inline tpb_t cores; the MLA index is this process's dense [0,32) device slot, used as the topology key by the collectives RID map. | mla_t (tdrv_ctx_0.mla[32]) | runtime/device-book.md | HIGH |
| NC (NeuronCore) | A physical NeuronCore / TPB — one of num_tpb per device (2 on V2, 8 on V3/V4), the five-engine machine (PE·ACT·POOL·DVE·SP + TopSP + Q7); physical_core_t keyed by device_tpb_idx. | physical_core_t (24 B) · device_tpb_idx @+0x04 | arch/overview.md | HIGH |
| NCFW | "Neuron Collective FirmWare" — the on-device Tensilica Xtensa-LX sequencer images (one IRAM + one DRAM blob per generation) that each NeuronCore's TPB sequencer runs to choreograph collective DMA; carried (not run) by libncfw.so. | libncfw_get_image @0x1179 · v{2,3,4,4_plus}_ncfw_{iram,dram}_bin | firmware/overview.md | HIGH |
| NDS | The Neuron DataStore — a process-shared, mmap-backed counter plane (one 67200-byte blob per (device, PID)) that publishes per-device/per-core atomic counters and model metadata for out-of-process observers; FNV-1a-hash-verified records + LOCK-atomic counters. | nds_open @0x5070d0 · libnds.a | datastore/overview.md | HIGH |
| NEFF | "Neuron Executable File Format" — the self-contained model package handed to nrt_load: a 1024-byte flat header + gzip-tar payload of JSON manifests and raw <eng>.bin blobs. No 4-byte magic; the loader keys off a UUID. | neff_header_t (1024 B) · neff_parse @0x4ca3f0 | neff/overview.md | HIGH |
| NQ (notification queue) | A circular buffer in host or device DRAM that hardware (compute engines / TopSP) produces completion/notify/error/trace messages into; userspace drains it by mmap. The kernel only allocates the ring and programs three registers (BASE_LO/HI, F_SIZE); a 0x28-byte register block per NQ instance. | NOTIFIC_NQ_SIZE 0x28 · nnq_init (neuron_nq.c) | kernel/notification-queues.md | HIGH |
| ntff | "Neuron Trace File Format" — the AWS-authored protobuf schema (ntff.proto + neuron_trace.proto, one package ntff) for all profiler output; 40 messages / 207 fields, recovered from the embedded FileDescriptorProto blobs. | descriptor_table_protodef_ntff_2eproto @0xad4200 | trace/ntff-format.md | HIGH |
| pring | "Physical ring" — a dma_ring_info_t (32 B) binding a dmem_t device-DRAM chunk to a descriptor count and a tx/rx role; a vring is dumped into a pring to commit a finished DMA program to hardware-visible memory. | dma_ring_info_t (32 B) · dma_ring_info_alloc | dma/virtual-rings.md · runtime/tdrv-dma-rings.md | HIGH |
| ptpb | "Physical TPB" — the physical NeuronCore slot (tpb_t, one of 8 inline per mla_t), as opposed to the virtual/LNC core threaded over it; addressed by (device_id, device_tpb_idx). | tpb_t (mla_t._tpbs[8]) | runtime/device-book.md | MEDIUM |
| RDH | "Reduction-Distribution-Hub" — a hub-centric reduce/distribute collective family (split intra/inter), a metaring_type (=3); reduces to a hub and distributes from it (star/hub, not a binary tree). | alg_inter_rdh_init @0xfe1f0 · enc/rdh.cc | collectives/hierarchical-rdh.md | HIGH |
| RID | "Routing-ID" — the fabric/NeuronLink routing address the collectives/switch layer addresses cores by; loaded at init from the driver's host-did→RID map (or a per-chip-family permutation fallback), distinct from the MLA index. | host_did_to_rid_map[32] · encd_get_rid @0x235000 | runtime/device-book.md | HIGH |
| S2M | Stream-to-memory — the UDMA Rx descriptor class (stream → host/device memory); the receive-side counterpart of M2S, with its own embedded-sem field pack. | aws_hal_sdma_s2m_set_embedded_sem_fields | dma/meta-ctrl-overlays.md | HIGH |
| SDMA | The per-arch KaenaHal engine band that programs the AnnapurnaLabs UDMA Stream-DMA engine for compute-while-copy and notification (transpose builder, notification CSRs, embedded-sem, MARIANA Q7/RDM/seed/FP8 slice). The engine band, not the controller. | aws_hal_sdma_* · al_hal_sdma.c | runtime/arch-sdma.md | HIGH |
| SP / TopSP | SP is the per-NeuronCore Sequencer/Scalar Processor (engine index 4 of PE/ACT/POOL/DVE/SP). TopSP ("Top Sync Processor") is a per-device orchestration engine (V2: 6, V3/V4: 16) that choreographs collectives; its instruction stream is DMA-fed, and it produces into NQs. | al_hal_tpb_eng_type SP=4 · struct ndhal_topsp | kernel/topsp.md | HIGH |
| STPB | The per-arch state/sync-TPB engine-init layer — the KaenaHal leaf bodies (aws_hal_stpb_*) that program the on-TPB compute engines (PE/ACT/POOL/DVE/SP) during bring-up: regs/ucode/DMA/run-stall, plus the Pooling DGE program. | aws_hal_stpb_{sunda,cayman,mariana}_* · aws_hal_stpb (1104 B) | runtime/arch-stpb.md | HIGH |
| SUNDA | Silicon codename for NeuronCore-v2 = Trainium1 + Inferentia2; 2 NeuronCores (two PCI device-IDs share one arch), arch enum value 2. | AL_HAL_TPB_ARCH_TYPE_SUNDA = 2 | arch/coretype-numbering.md | HIGH |
| TPB | The NeuronCore's compute tile — the five-engine machine PE·ACT·POOL·DVE·SP (+ TopSP, + GPSIMD/Q7); "TPB" and "NeuronCore" name the same physical core, counted by tdrv_arch_get_num_tpb_*. | al_hal_tpb_eng_type (PE=0…SP=4) | arch/overview.md | HIGH |
| UDMA | The AnnapurnaLabs/Alpine DMA controller (udma_q hardware queues + doorbell) that the SDMA band programs and the kernel owns the hardware ring/doorbell of; the 16-byte al_udma_desc is its wire descriptor. | al_udma_* · udma_q (kernel) | kernel/udma-main.md | HIGH |
| vNC | "Virtual NeuronCore" — the runtime's virtual-core abstraction threaded over one-or-more physical TPBs so a model can be placed on a subset/superset of cores (an LNC collapse); back-pointer in physical_core_t+0x08. | virtual_core_t (const, physical_core_t+0x08) | arch/overview.md | MEDIUM |
| vring | "Virtual ring" — the host-side staging structure (paged, growable, tx+rx doubly-linked list of 16-byte descriptors) the runtime authors a DMA program into and edits post-hoc, before dumping it into a pring. | vring_t (344 B) · vring_dump_to_pring_descriptors_padded @0x311f10 | dma/virtual-rings.md | HIGH |
| vtpb | "Virtual TPB" — the vNC's per-core information/pooling layer (tdrv/vtpb.c/vtpb_info.c) that resolves cross-subgraph remote variables and shares state-buffer init across the LNC-collapsed cores. | vtpb_info_shared_init_vtpb @0x3148e0 | runtime/tdrv-lifecycle.md | MEDIUM |
NOTE — the
enc/encdrows above intentionally point at the two distinct owning pages (collectives/engine-core.mdfor the composer,collectives/encd-overview.mdfor the emitter). The single-letter difference is the whole distinction; see the Confusable Pairs callout.
Confusable Pairs
These seven pairs share a name, a layer, or a direction and are wrong-by-default if transposed. Each callout gives the distinguishing symbol and the page that owns the disambiguation.
GOTCHA — NC vs vNC. NC (NeuronCore) is the physical core —
physical_core_t, keyed bydevice_tpb_idx, 2 on V2 / 8 on V3/V4, addressed by per-core register math. vNC is the virtual core —const virtual_core_t*atphysical_core_t+0x08, the LNC abstraction that collapses N physical cores into one logical unit (bounded byMAX_VIRTUAL_TPB). A model is placed on a vNC; geometry/CSR math runs on the NC. Routing per-core register addressing through the vNC, or counting vNCs where physical TPBs are meant, mis-addresses silicon. Physical: arch/overview.md →physical_core_t. Virtual: same page,physical_core_t+0x08; collapse bound in runtime/tdrv-lifecycle.md.
GOTCHA — SP/TopSP vs STPB. SP is the per-core Sequencer/Scalar Processor engine (
al_hal_tpb_eng_type SP=4); TopSP is the per-device "Top Sync Processor" orchestration engine (struct ndhal_topsp, 6 on V2 / 16 on V3/V4) — both are engines. STPB is not an engine at all: it is the per-arch state/sync-TPB engine-init layer (aws_hal_stpb_*leaf bodies) that programs the on-TPB engines at bring-up. A reimplementer who treats STPB as a hardware block, or who programs SP registers where the STPB init sequence belongs, mis-models the bring-up. Engines: kernel/topsp.md. Init layer: runtime/arch-stpb.md →aws_hal_stpb(1104 B).
GOTCHA — enc vs encd.
encis the host-side collective composer (enc/enc.cc,enc_context,enc_post_operation @0x11f790) — it decides ring order, mesh peering, reduce schedule.encdis the device-resident emitter (tdrv/encd.c,encd_context,encd_dma_mark_end @0x237200) — it materializes those decisions into SPAD op-streams and vring DMA packets. The edge is strictly one-directional:enccalls down intoencd;encdnever calls back. Wiring a libnccom/nccl*call into theencdfloor (or expectingencdto choose an algorithm) mis-places the boundary. Composer: collectives/engine-core.md. Emitter: collectives/encd-overview.md.
GOTCHA — kelf vs kbin.
kelfis the input — the KELF JSON descriptor inside the NEFF (kelf::kelf::load @0x497dc0) naming per-subgraphdef.json+<eng>.bin.kbinis the output — the 424-byte in-RAM device-binary descriptor recordkelf2kbinlowers each subgraph into (gen_kbin @0x4af930). There is no.kbinfile on disk; the NEFF carries the inputs to lowering, andkbin/model_tare the outputs that live only in RAM and device DRAM. Treating "KBIN" as an on-disk file symmetric with the JSON manifests is the canonical wrong assumption. Both owned by neff/kelf2kbin.md (lowering) and neff/kbin-structs.md (the record); the no-file fact is established in neff/overview.md.
GOTCHA — M2S vs S2M. M2S (memory-to-stream) is the Tx direction — memory → stream — packed by
al_sdma_m2s_build_*_descriptor. S2M (stream-to-memory) is the Rx direction — stream → memory — with its own embedded-sem field pack (aws_hal_sdma_s2m_set_embedded_sem_fields). The names are mirror-images and so is the data flow; transposing them sends a transfer the wrong way. Both owned by dma/meta-ctrl-overlays.md; the per-arch build leaves are in runtime/arch-sdma.md.
GOTCHA — SDMA vs UDMA. SDMA is the per-arch KaenaHal engine band (
aws_hal_sdma_*,al_hal_sdma.c) — the silicon-specific code that programs the DMA engine for compute-while-copy and notification. UDMA is the AnnapurnaLabs/Alpine DMA controller itself (al_udma_*, the hardwareudma_qqueues and doorbell) — the thing SDMA programs and the kernel owns the ring/doorbell of. The 16-byte descriptor is UDMA's wire format; SDMA is the band that builds and feeds it. A reimplementer conflating the programmer with the controller mis-partitions the HAL. SDMA band: runtime/arch-sdma.md. UDMA controller: kernel/udma-main.md.
GOTCHA — mesh vs metaring vs kangaring. These are three different collective algorithm families, not synonyms. mesh is the fully-peered
N·(N-1)family (alg_mesh_init @0x135990,enc_alg_mesh, device patternMESH=1) — the only family that servesALLTOALL. metaring is the ring-family schedule object and its topology sub-axis (enc_alg_metaring,metaring_type), shared by all ring-family algorithms via one composer. kangaring is one value ofmetaring_type(=1) — Neuron's multi-rail ring variant (alg_kangaring_init @0xfe480) — i.e. kangaring is a kind of metaring, while mesh is a separate family entirely. A reimplementer treating "metaring" as "ring" only, or "kangaring" as a peer of "mesh" rather than a ring-family member, mis-builds the composer dispatch (init_metaring_algorithm @0xfe570branches to ring/kangaring/rdh — never mesh). All three owned by collectives/algorithm-taxonomy.md.
Cross-References
- String Domain — the binary-embedded string corpus (env vars,
__FILE__paths, NEFF/kelf/kbin/ntff tokens) that surfaces most of this vocabulary - Coretype Numbering Reconciliation — the SUNDA/CAYMAN/MARIANA codenames and the six integer numbering schemes they appear under
- Subsystem ↔ Binary ↔ Source-TU Matrix — the term → source-TU → owning-page index this glossary cross-cuts
- The Five Binaries & Subsystem Map —
libnrt/libncfw/libnrtucode_extisa/libnccom/libndsand the four bind models the firmware/microcode terms hang off