Keyboard shortcuts

Press or to navigate between chapters

Press S or / to search in the book

Press ? to show this help

Press Esc to hide this help

Abbreviations & Symbol Index

This is the fast-lookup index for every acronym, opcode-mnemonic prefix, register name, firmware idiom, runtime symbol family, CSR/address block, and collective term that recurs across the wiki. It complements the Master Glossary: the glossary defines (a crisp paragraph + canonical anchor + deep page per term); this index cross-references (a one-line expansion + a link to the page that owns the term), and it adds the symbol-prefix families the glossary does not enumerate (nrt_*, nrtucode_*, aws_hal_q7_*, kmgr_*, tdrv_*, enc_*, ivp_*, module__xdref_*).

How to use it. Land here when you hit a bare token in another page — cc_op, tdrv_ctx, ivp_, P%i:, NQ, EVT_SEM — and want "what is it, and where is it defined". Find the row, read the one-line expansion, click the defining page. For the full definition + anchor, follow the glossary link in the row (or the Master Glossary directly). For the byte-level encoding lookup, see the Master ISA Encoding Appendix.

Provenance. Every expansion is derived from static analysis of the shipped binaries — recovered symbols, embedded prototype strings, config tokens, opcode-name string pools, and device .rodata literals. Symbol-family counts are reproducible: an nm/strings/rg against the one named binary by absolute path reproduces them. Where a row's meaning is inferred rather than read, it carries a confidence tag ([OBSERVED] / [INFERRED] / [MED]); most are [HIGH/OBSERVED] from symbols and are tagged only where confusable. The five NOTE / GOTCHA callouts at the end disambiguate the genuinely-confusable abbreviations (the v5-label overload, SP vs TOP_SP, ct37 vs arch_id 36, XRP the not-Cadence-framework, CCE the not-an-engine).

Literal | inside a cell is escaped \|. A row that points at the glossary means the term is defined there; a row that points at a subsystem page means that page owns the deep treatment.


1. Engine & codename abbreviations

The generation codenames, the coretype/arch_id identity bytes, and the engine ordinals. Deep table: Codename ↔ Generation Cross-Walk and the Codename Cross-Walk Table.

termexpansion (terse)defining page
GPSIMDGeneral-Purpose SIMD substrate = the POOL engine, 8× Vision-Q7 Cairo DSP coresWhat GPSIMD Is
CairoTensilica µarch name of the GPSIMD core (uarchName = Cairo); one frozen Vision-Q7 NX configglossary · Cairo
Vision-Q7 / Q7the off-the-shelf Cadence Tensilica vector DSP IP that GPSIMD isCore Identity & Configuration
ncore2gpTensilica CoreID of the Cairo config; the XTENSA_CORE= handle for the device disassemblerCore Identity & Configuration
coretypefirmware byte identifying a generation — {6,13,21,29,37} [HIGH/OBSERVED]Codename Cross-Walk
arch_idhost NCFW image-selector byte; arch_id = coretype − 1 (the only uniform relation — the coretype set {6,13,21,29,37} is not a flat +8 stride). {5,12,20,28} OBSERVED (double-anchored v2–v4); 36* INFERRED (v5, no NCFW byte)Codename Cross-Walk
SUNDAgen v2 (NC-v2; Trn1/Inf2) · coretype 6 / arch_id 0x05 · the no-HW-Decode floorSUNDA v2 Baseline
CAYMANgen v3 (NC-v3; Trn2) · coretype 13 / arch_id 0x0c · the byte-grounded reference genCodename ↔ Generation Map
MARIANAgen v4 (NC-v4) · coretype 21 / arch_id 0x14 · distinct compile of the CAYMAN contractCodename ↔ Generation Map
MARIANA_PLUSgen v4+ (Trn3-pre) · coretype 29 / arch_id 0x1c · feature-flag delta, byte-identical EXTISAMARIANA_PLUS Delta
MAVERICKgen v5 (NC-v5) · coretype 37 OBSERVED / arch_id 36 INFERRED · header-only, no NCFW imageMAVERICK Profile
TONGAlegacy NC-v1 (Inf1) outlier — not a 6th gen; no coretype/arch_id/NCFWCross-Gen Opcode Diff + TONGA
NC-v1 … NC-v5the NeuronCore-version axis (v1=TONGA … v5=MAVERICK)Codename Cross-Walk
EXTISAper-gen statically-embedded Vision-Q7 ext-ISA device image (ELF32-Xtensa, e_machine=94)EXTISA Q7 SO-Blob Inventory
TPBTensor-Processing Block = one NeuronCore as a block of engines on one SBUFThe Seven Faces
PEProcessing-Element array (128×128 systolic Tensor engine) · engine_idx 0 · sole PSUM writerPE Matrix-Multiply Path
ACTScalar / activation / PWL engine · engine_idx 1 · folds into DVE on MAVERICKActivate + PWL
POOLthe GPSIMD engine (8× Vision-Q7 Cairo) · engine_idx 2 · subject of this wikiPOOL Dispatch Loop
DVEthe Vector engine · engine_idx 3 · predicated-op family; absorbs ACT on MAVERICKDVE State Read-Back
SP (TPB_SP)per-NeuronCore Sync/control front-end executor · engine_idx 4Per-Engine Firmware Depth
TOP_SPstandalone NX-core sequencer that walks the cc_op collective program · engine_idx 5TOP_SP Collective Lowering
NCFWNeuronCore Firmware — a separate scalar Xtensa-LX collective-management coreThe NCFW Scalar-LX Core
engine_idxfirmware engine ordinal (PE 0 … TOP_SP 5) — not the NKI compiler engine enumThe Compiler Map

2. ISA mnemonic families & opcode prefixes

The Vision-Q7 mnemonic family, the scalar base, and the recurring opcode bytes. Byte-level encodings: Master ISA Encoding Appendix.

prefix / opcodeexpansion (terse)defining page
ivp_*Vision-Q7 vector-instruction mnemonic family (IVP SIMD TIE package xt_ivp32); 1065 distinct mnemonics in the libisa-core.so opcode string poolISA Template & Partition
(scalar base)base-Xtensa LX7 scalar ops (addi, l32i.n, entry, retw.n) — not ivp_*-prefixedEight Register Files
0x41 TENSOR_TENSOR_ARITH_OPthe int32/uint32 add/sub/mul datapath (GpSimd-native lane)Collective + cc_op Enums
0xBF SB2SB_COLLECTIVEone SBUF→SBUF collective hop (a ring all-reduce step)S3D3 Collective (SB2SB)
0xF0 EXTENDED_INSTcustom-op extended-instruction space; DMAs its own Q7 EXTISA imagePOOL Extended-Opcode (0xF0)
0xC7/0xC8/0xD9PSEUDO_TRIGGER_ALL_REDUCE / _COLLECTIVE / _COLLECTIVE2 triggersCollective + cc_op Enums
0xCB/0xD5/0xD8PSEUDO_SEND_RECV / _SYNC_BARRIER / _CORE_BARRIERCollective + cc_op Enums
0xC3 PSEUDO_DMABARRIERDMA barrier opcodeCollective + cc_op Enums
0xDB PSEUDO_CUR_PROCESSING_RANK_IDread this core's PRID/rankPseudoCurProcessingRankID
NEURON_ISA_TPB_OPCODE_*the arch-ISA opcode enum (the byte→name table) in aws_neuron_isa_tpb_common.hCollective + cc_op Enums
TIETensilica Instruction Extension — the DB/language defining the Vision-Q7 coproc ISATIE Database

3. FLIX encoding terms

The Tensilica VLIW scheme. Deep decoder: FLIX Bundle-Decoding Methodology; byte tables: FLIX Encoding.

termexpansion (terse)defining page
FLIXFlexible-Length Instruction eXtensions — the VLIW scheme; 14 formats, 46 slotsFLIX Encoding
bundleone fetched FLIX word (16 B wide / 8 B narrow / 2–3 B scalar fallback)FLIX Encoding
formatone of the 14 bundle layouts (selector encoding)FLIX Encoding
slotone operation position within a format; 14 formats → 46 slotsFLIX Encoding
x24 / x16a / x16bthe three scalar formats (24-bit + two 16-bit-density)FLIX Encoding
F0/F1/F2/F3/F4/F6/F7/F11the eight wide (16-byte / 128-bit) bundle formatsFLIX Encoding
N0/N1/N2the three narrow (8-byte) bundle formatsFLIX Encoding
wide / narrowwide = 16 B / 4–5 slots (op0 E, or F odd b3lo); narrow = 8 B / 2–4 slots (F even b3lo)FLIX Encoding
length_table256-entry table → "7 length-class outcomes → 4 byte-lengths {2,3,8,16}"FLIX Decoding
FLIX-desynca linear byte sweep desyncs on literal pools / the F0 8-vs-16 split → bundle interiorFLIX Decoding
b3lolow nibble of byte 3 — the wide-vs-narrow + 8-vs-16 discriminator for op0==0xFFLIX Decoding

4. Register-file names & ISA metadata fields

The eight register files (two scalar, six Vision-Q7 SIMD) and the per-operand metadata tables. Deep page: The Eight Register Files.

name (short)expansion (terse)defining page
AR (a)scalar address/general registers (32-bit × 64, windowed ABI)Eight Register Files
BR (b)scalar boolean registers (1-bit × 16; BR2/4/8/16 views)Eight Register Files
vec (v)512-bit SIMD vector file (32 × 16-bit lanes; 32 regs)Eight Register Files
vbool (vb)per-lane SIMD predicate (boolean) file (64-bit × 16)Eight Register Files
valign (u)alignment registers priming unaligned vector loads (512-bit × 4)B06 Vector Loads + valign
wvec (wv)1536-bit wide MAC accumulators (quad-width readout; 4 regs)Eight Register Files
b32_pr (pr)64-bit predicate/pack registers (int64pr ctype; 16 regs)Eight Register Files
gvr (gr)global/state vector registers (gsr ctype, flags 0x0d; 8 regs)Eight Register Files
SuperGatherthe IVP two-phase gather-register-staged gather/scatter unitB19 SuperGather
ctypeTIE C value-type a register/operand maps to (64-entry ctypes table)ctype/coproc/funcUnit Tables
coprocthe coprocessor (exactly one: {name="Vision", number=1})ctype/coproc/funcUnit Tables
funcUnitfunctional unit an opcode binds (one: XT_LOADSTORE_UNIT, 2 copies)ctype/coproc/funcUnit Tables
MSXEA3 dispatch-mode/state register (SR 0xe5) — software-managed dispatch demuxXEA3 Interrupt Architecture
VECBASEvector-base SR 0xe7 — the [31:6] page all dispatch vectors hang offXEA3 Interrupt Architecture
PRIDspecial-register 235 — the SPMD rank of a Q7 core (PRID_ID & 0xF ∈ {0..7})The 8-Core SPMD Model

5. Firmware & device symbol idioms

The device-firmware build idioms used as recovery anchors. The trace tags are confirmed in the device .rodata of libnrtucode_internal.so (S: literals 16×, P%i: literals 1016×).

idiomexpansion (terse)defining page
kernel_info_tableper-image opcode→handler map; 8-byte entries {0,0,spec,opcode,u32 funcVA} (CAYMAN: 17)kernel_info_table Layout
.xt.propper-function FLIX property sections (.xt.prop.<mangled>) — the FLIX-aware code-mode anchor (138× in libnrtucode_internal.so)FLIX Decoding
'S:' tagDEBUG trace string of the SEQ control-engine stream (e.g. S: Dispatch opcode=0x%x, S: BEGIN on cayman)SEQ Decode / Dispatch Hub
'P%i:' tagDEBUG trace string of the per-CPU POOL/Q7 data-plane stream; %i = the core's PRID (e.g. P%i: SB2SB_Collective ...)SEQ Decode / Dispatch Hub
HW-Decodethe later-gen hardware opcode-decode CAM dispatch mode (SUNDA disables it)HW-Decode vs Sunda Dual Fetch
Sunda modethe SUNDA software dual-fetch path (S: NX in Sunda mode: HW decode disabled)HW-Decode vs Sunda Dual Fetch
SEQthe device sequencer front-end (boot, FSM, fetch + PC-redirect, decode/dispatch hub)SEQ Main FSM Loop
DGEDescriptor Generation Engine — 3-backend selector (Pool / RTL / software)DGE 3-Backend Selector
Boot handshakeunbooted sentinel 0x6099CB34 → host CAS-writes claim 0x502B2DA1Boot / Reset Sequence
XCHAL_*the Tensilica config-token prefix (XCHAL_VISION_TYPE, XCHAL_XEA_VERSION, …) in core-isa.hCore Identity & Configuration

6. Runtime / host symbol families

The host x86-64 runtime prefixes. Counts are reproducible with nm/strings against the one named binary; the host families (aws_hal_*, kmgr_*, tdrv_*, enc_*) live in libnrt.so as embedded prototype strings (the library is not stripped — these are its symbolized C++ prototypes). nrtucode_* is a real .symtab family in libnrtucode_internal.so.

prefixexpansion (terse)binary + countdefining page
nrt_*host-runtime public API (nrt_load, nrt_set_pool_eng_ucode, …)libnrt.so · 121 dynsymlibnrt Surface Map
nrtucode_*host micro-code subsystem managing Q7 images (resolvers, ll-load/unload, opset)libnrtucode_internal.so · 60 nmnrtucode Subsystem
aws_hal_q7_*host Q7 HAL — register/window/swap-table accessors (aws_hal_q7_swap_table, …)libnrt.so · 99 stringsThe aws_hal_q7_* HAL
aws_hal_*the wider KaenaHal hardware-abstraction surfacelibnrt.so · 2528 stringsThe aws_hal_q7_* HAL
kmgr_*host kernel/model manager (kmgr_init, kmgr_sync_exec, kmgr_exec_resources_t)libnrt.so · 471 stringsExec-State Census
tdrv_*host tensor-driver layer (tdrv_ctx_t, tdrv_arch_get_num_tpb, …)libnrt.so · 1226 stringsExec-State Census
tdrv_ctx_tthe process-global host root struct (tdrv_ctx_0, embeds mla[32])libnrt.so (DWARF)Exec-State Census
enc_* / encd_*host-side collective-communication (CC) program builder family (enc_alg_type, enc_comm_info, encd_context)libnrt.so · 10208 strings (ENC_ALG/ALLGATHER 38)Device-Firmware Globals §1.5
module__xdref_*the libfiss-base.so per-lane value-leaf family (module__xdref_add_16_16_16)libfiss-base.so · 864 nmfiss 864-Leaf Value Oracle
ncfw_*NCFW host logging/scheduler symbols (ncfw_log_spad_ctrl_cc_op_entry)libncfw.so · 156 nmNCFW spad-ctrl cc_op
nrtucode_core_tper-core host handle (0x70 B; coretype byte at +0x10)libnrtucode_internal.sonrtucode_core_t Struct
nrtucode_context_tper-model host context (0x28 B; lifecycle + dmem allocator)libnrtucode_internal.sonrtucode_context_t
ISSInstruction-Set Simulator (host config libs used as a live value oracle)libfiss-base / libcas-coreISS Semantic Synthesis
xdrefthe value-leaf naming stem (<op>_<wout>_<win0>_<win1>) — "cross-decode reference value fn"libfiss-base.sofiss Value Oracle

7. CSR & address-block names

The device control/status-register block names and SoC address regions. Block-name index: CSR Field-Table Index.

block / regionexpansion (terse)defining page
udma_m2sal_udma memory-2-stream outbound descriptor engine CSR blockCSR — udma_m2s
udma_s2mal_udma stream-2-memory inbound descriptor engine CSR blockCSR — udma_s2m
udma_gen / udma_gen_ex / tdma_modelshared SDMA control CSR blocksCSR — udma_gen/tdma
al_udmathe Annapurna-Labs micro-DMA hardware engine the udma_* blocks programal_udma HW DMA Engine
NOTIFIC / notific_n_queuehardware instruction-notification queue → coalesces into SW NQ ringsCSR — NOTIFIC Queue
EVT_SEMevent-semaphore CSR region (*_EVT_SEM_SEMAPHORE_SET/READ/DEC_BASE; 2649 strings in libnrt.so)Address — EVT_SEM Regions
FIS / fis_controlper-fabric-master error-trigger / isolation block (acronym not spelled in corpus)CSR — FIS + errtrig + spad
sprotthe security-protection (firewall) vector within the FIS/error-routing fabricCSR — FIS + errtrig + spad
errtrigthe error-trigger vector routing faults (e.g. APB timeout) into the interrupt fabricCSR — FIS + errtrig + spad
intc_*the interrupt-controller CSR blocks (1-group APINTC / 4-group)CSR Field-Table Index
SBUF / STATE_BUFon-chip State Buffer — shared 32 MiB @ SoC 0x2000000000 (Q7-reachable)SBUF + PSUM Bank Model
PSUMPE-private accumulator bank — structurally unreachable from GPSIMD (no AXI aperture)SBUF + PSUM Bank Model
SPADthe scratchpad the NCFW / TOP_SP walks the cc_op program inNCFW spad-ctrl cc_op

8. Collective & NCFW vocabulary

The collective-communication control plane: the program the NCFW/TOP_SP walks, the host-side encoder family, the messaging transport, and the DMA datapaths.

termexpansion (terse)defining page
cc_opthe collective-communication operation program the NCFW/TOP_SP walks (per-step spad-ctrl descriptors); "cc_op" string @ 0x650d4 in libncfw.soNCFW spad-ctrl cc_op + tsync
cc_op_infothe per-entry descriptor record ({cc_op, op_type, alg, alg_name})NCFW spad-ctrl cc_op
tsyncthe time-sync sequencing layer realising a collective step-by-stepNCFW spad-ctrl cc_op
enc_op_typethe high-level op enum the op-list lowers (ENC_ALLGATHER=0ENC_ALLTOALL_V=12, 13 kinds)Device-Firmware Globals §1.5
enc_alg_typethe algorithm enum (ENC_ALG_RING=0, _HIER=1, _MESH=2, _KANGARING=3, … _INVALID=11)Device-Firmware Globals §1.5
enc_alg_mesh_typethe mesh sub-selector (FULL_MESH=0INVALID=4)Device-Firmware Globals §1.5
enc_comm_typetopology scope (H_COMM_INTRA_ID=0, H_COMM_INTER_ID=1)Device-Firmware Globals §1.5
enc_comm_infothe 72-byte communicator topology descriptorDevice-Firmware Globals §1.5
EVT_SEMevent-semaphore — the engine-quad {INST_START, INST_END, EXPLICIT, EVT_SEM} notify member + the CSR regionAddress — EVT_SEM Regions
NQSW Notification Queue — up to ten software-owned ring queues the NOTIFIC block writes over AXICSR — NOTIFIC Queue
CCECompute-DMA / in-transfer compute — reduction inside the SDMA transfer (CDMA = DDMA + CCE)CCE In-Transfer Compute
SB2SB"State-Buffer to State-Buffer" — the 0xBF intra-/inter-die collective hopS3D3 Collective (SB2SB)
RDMAcross-core/cross-die SBUF→SBUF peer-to-peer byte movement over the die meshRDMA Cross-Die P2P
XRPhost↔DSP collective-control messaging transport — a bespoke Annapurna queue, not Cadence XRPXRP Host↔DSP Messaging
pring / vringpersistent vs template DMA-descriptor rings (NCFW copies vringpring once, reuses)pring Descriptors
ncfw_log_spad_ctrl_cc_op_entrythe NCFW logger that emits a cc_op spad-ctrl entry (the cc_op string anchor)NCFW spad-ctrl cc_op

9. Container, compiler & general acronyms

termexpansion (terse)defining page
NEFFNeuron Executable File Format — 1024-byte header + inner gzip→tar containerNEFF Container Byte Format
metaneffthe NEFF host-side protobuf I/O key-ring (MetaNeff/MetaTensor; binds tensor→var_idmem_ref)metaneff I/O ABI
BIR / Penguinthe neuronx-cc "penguin" backend IR (libBIR.so, ~110 Inst* classes)Penguin BIR Instruction Set
NKINeuron Kernel Interface — the Python frontend (engine enum tensor=1…sync=6, not engine_idx)NKI Frontend
MXmicroscaling — the block-scaled low-precision dequant/compute path (MXTENSOR_*)MX Microscaling Path
SundaISelthe compiler instruction-selection pass for the GPSIMD/Sunda lane (PSUM fall-back enforcer)SundaISel Deep-Dive
CoreV5 / core_v5compiler ArchLevel slot = MARIANA_PLUS / Trn3-PRENOT MAVERICK (see §10 NOTE)Codename Cross-Walk
ISS / FISS / CASInstruction-Set Sim / Fast ISS (value) / Cycle-Accurate Sim (timing, license-gated)ISS Semantic Synthesis
XEA3Xtensa Exception Architecture 3 — single unified DispatchVector, MS-managed dispatchXEA3 Interrupt Architecture
Xtensa24the core's ISA name (arch = Xtensa24) — the LX7 family (windowed ABI + 16-bit density + FLIX)glossary · Xtensa24/XEA3
RTTI absencedevice firmware is -fno-rtti_ZTS/_ZTI/_ZTV count is 0 across all 29 device ELFsMethodology
OBSERVED / INFERRED / CARRIEDprovenance tags (read-from-artifact / reasoned-over / reused-from-prior)Confidence & Walls Model
HIGH / MED / LOWconfidence tags (trust level), orthogonal to provenanceConfidence & Walls Model
Walla genuine static-analysis boundary (closable-with-license / -corpus / -hardware / fundamental)Confidence & Walls Model
QUIRK / GOTCHA / CORRECTION / NOTEthe inline callout markers used on every pageHow to Read This Guide

10. Disambiguation callouts — the confusable abbreviations

GOTCHA — the "v5" label is overloaded three ways; none derives the other. CoreV5 / core_v5 is the compiler ArchLevel slot (host platform / libwalrus), and it maps to MARIANA_PLUS / Trn3-PREcoretype 29 / arch_id 0x1c, an NCFW-present gen. NC-v5 / GPSIMD coretype 37 is the genuine 5th siliconMAVERICK — with no NCFW image. NeuronCoreVersion::V5 = 5 is the arch-ISA cap enum (present in maverick, absent in mariana). A page citing "v5" must say which axis. See Codename Cross-Walk §CoreV5-vs-NC-v5 and MAVERICK Profile §1.

GOTCHA — SP vs TOP_SP are two distinct engines. SP (TPB_SP) is the per-NeuronCore Sync/control front-end executor (engine_idx 4); TOP_SP is the standalone NX-core sequencer that walks the cc_op collective program in SPAD (engine_idx 5). They share the _SP suffix and the TPB_SP token but are not the same core. See TOP_SP Lowering.

NOTE — ct37 is OBSERVED, arch_id 36 is INFERRED. MAVERICK's coretype 37 is anchored twice in firmware (the header enum ordinal and the twin resolver bitmasks) → [HIGH/OBSERVED]. Its arch_id = 36 (0x24) has no cmp $0x24 anywhere and no v5 NCFW image — it is the coretype − 1 extrapolation only → [MED/INFERRED]. Never hard-code arch_id 36. See Codename Cross-Walk.

NOTE — XRP is not a Cadence framework here. Despite the name there is no Cadence XRP present in the corpus; the host↔DSP transport is a bespoke Annapurna-Labs message queue. The only xrp substring corpus-wide is libnrt's unrelated AF_RXRPC. See XRP Host↔DSP Messaging.

NOTE — CCE is a datapath, not an engine, and NCFW is a core, not the FLIX one. CCE (Compute-DMA) is the in-transfer reduce path inside the SDMA fabric — there is no "CCE engine". NCFW is a separate scalar Xtensa-LX management core: do not run the FLIX decoder on its images (the spurious "~26–28% FLIX" artifact comes from doing exactly that). Three cores, three ISAs: Q7 (Vision-Q7 FLIX, this wiki), NCFW (scalar-LX), TOP_SP (NX sequencer). See Keystone Facts and The NCFW Scalar-LX Core.


See also