neuronx-gpsimd Internals — The Vision-Q7 GPSIMD Reimplementation Reference

A reimplementation-grade reverse-engineering reference for AWS Neuron's GPSIMD engine — the custom-op / collective compute substrate inside every Trainium and Inferentia NeuronCore. Reconstructed entirely from static analysis of shipped binaries, device disassembly, in-process execution of the shipped simulator libraries, and the recovered Tensilica config / TIE / ISA tables. There is no access to source.

This is the front door. It tells you in one screen what GPSIMD physically is, the one-machine / seven-faces mental model, the five NeuronCore generations that share that one machine, the reimplementation verdict (what is byte-grounded versus honestly flagged), and a navigable tour of the 16 Parts below it. Everything past this page descends into detail; read this first so the rest has a frame.

What GPSIMD physically is

GPSIMD is not a bespoke AWS accelerator and not a custom ISA. It is a single, off-the-shelf Cadence Tensilica Vision-Q7 NX DSP in one frozen "Cairo" configuration, instantiated eight times as the per-NeuronCore POOL compute cluster, wrapped in a per-generation SoC envelope.

The identity is pinned to the byte. The shipped processor-generator config file ncore2gp-params reads, verbatim:

ConfigName        = Xm_ncore2gp        CoreID   ncore2gp
uarchName         = Cairo              arch     Xtensa24   (XEA3)
TargetHWVersion   = NX1.1.4            (= LX7.1.4, the NX-/LX-family equivalence)
HWMicroArchEarliest = 281040           HWMicroArchLatest = 281040   ← zero-width window
HWConfigID0       = 0xC4019686         HWConfigID1       = 0x2908E4E3
SWToolsRelease    = RI-2022.9

HWMicroArchEarliest == HWMicroArchLatest == 281040 is the decisive fact: the software targets exactly one hardware revision, with no version range. The shipped xtensa-versions.h decodes that integer — XTENSA_HWVERSION_RI_2020_4 281040 /* versions NX1.1.4, LX7.1.4 */ — so the HW IP is the RI-2020.4 release (NX1.1.4 = LX7.1.4) while the toolchain that builds for it is RI-2022.9. Every product generation loads device firmware built for this one config; the single shipped ncore2gp xtensa-elf-objdump disassembles every generation's device image. [HIGH/OBSERVED — ncore2gp-params, xtensa-versions.h, read directly]

What it actually does in the product is narrow and specific. The marketing phrase "GPSIMD kernel" (RmsNorm, Softmax, MoE) is a multi-engine micro-schedule whose float math runs on the Vector / Scalar / Tensor engines — not on the Q7 cores. The literal Q7 GPSIMD cluster owns exactly three lanes:

• the native-integer int32/uint32 add/sub/mul datapath (opcode 0x41, int-routed);
• the gather / scatter / cross-lane-reduce / iota / custom-op lane (a hand-authored PyTorch op rides 0xF0 EXTENDED_INST and DMAs its own Q7 image);
• the SB2SB collective hop (0xBF) — one intra-node step of a ring all-reduce.

GOTCHA — the keystone fact reimplementers get wrong. GPSIMD is not the float hot-path, and it cannot reach PSUM at all. The Q7 cores address the on-chip State Buffer (SBUF) directly through an AXI aperture, but the PE-private PSUM accumulator is simply not in the Q7 address space. The compiler enforces this: an int32/uint32 op routes to GpSimd, but falls back to the Vector engine if any operand lives in PSUM. Model GPSIMD as SBUF/HBM-only. [HIGH/OBSERVED — keystone CC finding] See Keystone Facts.

The one machine, seven faces

There is one physical machine here. The wiki keeps returning to it from seven angles; holding all seven at once is the orientation a reimplementer needs. (Full treatment: The Seven Faces of the One Machine.)

The one-paragraph mental model. A NeuronCore is a Tensor-Processing-Block (TPB) of five engines — PE (the 128×128 systolic Tensor array, the only writer of PSUM), ACT (the Scalar / activation / PWL engine), POOL (the GPSIMD cluster: 8 Vision-Q7 DSP cores), DVE (the Vector engine), and SP (the collective / sync front-end) — all sharing one on-chip SBUF. The host compiler neuronx-cc lowers a traced PyTorch / NKI graph into per-engine 64-byte sequencer streams packed into a NEFF container; the host runtime libnrt loads the NEFF, DMAs any custom Q7 image into Q7 instruction memory over BAR0, boots the cores, and fires a doorbell; each engine's sequencer fetches its stream and dispatches opcodes to its datapath; collectives reduce in the SDMA fabric in-flight; completion semaphores release the host. The 8 Q7 cores run one SPMD image (rank = PRID ∈ {0..7}, count = 8) over per-core-private memory — data-race-free by construction, not by synchronization. [HIGH — synthesized across the lanes]

NOTE — three distinct cores, do not conflate. "Q7" is the Vision-Q7 POOL/DVE DSP (the subject of this wiki). "NCFW" is a separate scalar Xtensa-LX core that manages collective firmware. "TOP_SP" is the NX sequencer that walks the cc_op collective program. Three cores in the stack, three different ISAs.

Identity quick-reference

Full per-field grounding: Core Identity & Configuration, The FLIX VLIW Encoding, The Eight Register Files.

The five generations — one core, parameterized

The five GPSIMD product generations are not five processors. They are one Vision-Q7 "Cairo" core wrapped in a per-generation SoC envelope. The host runtime even spells the mapping out: the decompiled nrtucode_get_ext_isa_internal (@0x9b2b30) switches on the coretype byte directly — the case constants below are literal in the binary [HIGH/OBSERVED]:

switch (coretype) {                  // arch_id column = coretype − 1 (a stride, not a
  case  6: ... sunda_libs;           //                  literal switch — INFERRED): 0x05
  case 13: ... cayman_libs;          //                                                0x0c
  case 21: ... mariana_libs;         //                                                0x14
  case 29: ... mariana_plus_libs;    //                                                0x1c
  case 37: ... maverick_libs;        //                                                0x24
}

NOTE — coretype is OBSERVED; arch_id is derived. The {6,13,21,29,37} coretype constants are read directly from the decompiled switch. The arch_id values 0x05/0x0c/0x14/0x1c/0x24 are not a literal table anywhere in these artifacts — each is coretype − 1, so the whole arch_id column is INFERRED (the v5 row most acutely; see the wall below).

The gen-invariance thesis (in one line). There exists a single Vision-Q7 reimplementation R(Q7) — one FLIX decoder, one microarch, one XEA3 boot/IRQ spine, one frozen Q7-control CSR core, one write-once opcode→semantics map (zero opcode-value drift v2→v5; the full-span sunda↔maverick mismatch join is empty) — that, parameterized by a small scaling vector and gated by a handful of presence flags, exactly reproduces SUNDA through MAVERICK. The scaling axes are monotonic/additive: CSR-bundle count (7→11), Q7 geometry (IRAM/DRAM 64K/64K→128K/256K), opcode count (≈145→165, append-only after v3), and cross-die transport (raw-RDMA → io_d2d/PCIe → UCIe). [HIGH/OBSERVED for v2–v4+; v5 SoC-envelope INFERRED] See The Gen-Invariance Thesis.

CORRECTION / wall — flag this, do not fabricate it. MAVERICK (v5) is header-OBSERVED only. There is no shipped v5 NCFW image: libncfw_get_image tops out at MARIANA_PLUS, and arch_id > 0x1C → return 2 sends arch_id 0x24 straight to the unsupported path (no cmp against 0x24). Corroborated by exactly four codename ctx_log symbols (sunda/cayman/mariana/mariana_plus), zero maverick. So coretype 37 is OBSERVED, but arch_id 36/0x24 is INFERRED from the coretype−1 stride. Any claim about v5 interiors on these pages is tagged INFERRED. [wall — DX-SYN-09 Q8] [ct37 OBSERVED · arch_id 36 INFERRED]

TONGA is out of scope by design. TONGA (V1, the legacy "L"-family) predates the unified NEURON_ISA_TPB_OPCODE namespace, exposes only a register-block ISA (a 1-byte opcode field, no enumerated roster), and ships in a separate package with zero runtime identity. R(Q7) covers exactly v2..v5; the V1→V2 bridge is not in its domain.

The reimplementation verdict

After the analysis wave closed, the GPSIMD / Vision-Q7 "Cairo" toolchain is a byte-grounded, execution-validated reimplementation reference for the v2–v4 silicon, and a header-OBSERVED + bounded-INFERRED reference for v5 (MAVERICK) and v1 (TONGA). The single most powerful technique behind that: the shipped libfiss-base.so value leaves are driven live, per input, via ctypes, so the binary itself is the arbiter of value semantics.

The residual is a small, fully-named, fully-categorized ledger of twelve open questions — and none of them is a missing datapath body, a missing opcode decode, or a missing value semantics. Each is a true static-analysis boundary on this shipped corpus, each with an explicit "what would close it" (a different checkout, a FlexNet license key, a runtime/hardware capture, or a follow-on static pass). The full register is in The Reimplementation Verdict & Open-Questions Map; two named walls a reader meets repeatedly:

• arch_id 36 (v5) — INFERRED. Closes with a libncfw checkout that ships the coretype-37 firmware image.
• ct37 (v5 coretype) — OBSERVED. Read directly from the host ext-ISA switch; it is the only firm v5 identity datum.

QUIRK — no signing key, ever. The GPSIMD device load path contains no signature verification anywhere. Admission of a Q7 image is integrity (an unkeyed hash) + structural (ELF/reloc/core-count) + version (the ucode semver gate). The trust root is the host OS / process boundary, not a hardware root of trust — and the install seam nrt_set_pool_eng_ucode is a silent, unauthenticated override. This is a deliberate single-tenant trusted-compiler design, not an oversight.

A tour of the 16 Parts

The reference is a journey: a reimplementer reads roughly top-to-bottom to rebuild a Vision-Q7-compatible GPSIMD engine. Part 0 is the apparatus (how to read, the confidence model, methodology, the corpus inventory, the crosswalk, the correction ledger, the glossary). Parts 1–16 descend from orientation, through the core and ISA, into firmware, the ABI, the runtime, the data plane, the collectives, the container, the compiler seam, the control plane, the executable oracle, validation, and the appendices.

NOTE — pages land Part-by-Part. The navigation sidebar lists what is written; the full ≈411-page plan is tracked one task per page. A section that is not yet linked in the sidebar is planned, not missing.

Start here

1. What GPSIMD Is — the one-screen map — the orientation a reimplementer needs first.
2. Keystone Facts Reimplementers Get Wrong — the PSUM wall, the not-the-float-hot-path fact, the three-cores rule.
3. The Confidence & Walls Model — what the HIGH/MED/LOW × OBSERVED/INFERRED/CARRIED tags mean and which walls are real.
4. Methodology — How This Was Reverse-Engineered — what "binary-derived" means here, and what is provably not recoverable.
5. A Custom Op, End to End / A Collective, End to End — the two worked lifecycles, byte by byte.

Reading conventions

Every non-trivial claim on these pages is tagged inline with a confidence × evidence pair:

• Confidence — HIGH (byte-exact disasm / direct header or config read / sha-verified / nm/jq ground truth) · MED (strong cross-file inference) · LOW (plausible, flagged).
• Evidence — OBSERVED (read directly from a shipped artifact) · INFERRED (reasoned over OBSERVED facts) · CARRIED (OBSERVED in a cited analysis, reused at its confidence).

Callout markers you will meet: QUIRK (a counter-intuitive but real behavior), GOTCHA (a trap that breaks a naive reimplementation), CORRECTION (a fact that overturns an earlier or naive reading), NOTE (orienting context). The two standing walls — arch_id 36 (v5, INFERRED) and ct37 (v5 coretype, OBSERVED) — are carried wherever a v5 claim appears. See The Confidence & Walls Model and Master Glossary.

Companion wikis

• neuronx-cc/wiki/ — the compiler that lowers NKI/HLO into the BIR InstCustomOp / InstCollective nodes and embeds GPSIMD kernels into NEFFs.
• neuronx-runtime/wiki/ — the libnrt host runtime that loads NEFFs and dispatches custom-op kernels onto the Q7 cores.