Neuron Runtime Internals

Reverse-engineering reference for the AWS Neuron runtime stack. Source packages: aws-neuronx-dkms_2.27.4.0 (unstripped GPL-2.0 kernel source) · aws-neuronx-runtime-lib_2.31.24.0-0b044f4ce (libnrt.so + libnds.a with DWARF; libncfw.so + libnrtucode_extisa.so unstripped but symtab-only, no DWARF) · aws-neuronx-collectives_2.31.24.0-1a31ba186 (libnccom.so with DWARF; libnccom-net.so symtab-only). Every address on a page pins to these builds.

What this wiki is

This book documents the runtime side of the AWS Neuron stack at reimplementation grade: the layer that takes a compiled model (a NEFF, produced by neuronx-cc), loads it onto Trainium/Inferentia silicon, runs inference, and orchestrates multi-device collectives. It spans five binaries — a userspace runtime, a GPL kernel driver, two on-device firmware/microcode carriers, and a forked NCCL — reconstructed function-by-function from static analysis.

The bar is the house bar: a competent systems engineer should be able to rebuild each component from its page — the algorithm, the data layout, the wire format, the decision logic — and tell which parts are certain and which are inferred. See How to Read This Book, which also covers the evidence and confidence conventions.

The five binaries

libnds.a (the Neuron DataStore) is statically linked into libnrt.so and also shipped standalone.

The layered stack

  CONSUMERS    libtorchneuron.so / libneuronpjrt.so / JAX-PJRT
                       │  NRT_2.0.0 + NRT_3.0.0 ABI (149 exports)
                       ▼
  RUNTIME      libnrt.so  ── load NEFF, plan memory, submit, harvest, trace
                       │            │ dlopen
                       │            ├─► libnccom.so ── multi-node collectives (EFA)
                       │            └─► libnrtucode_extisa.so ── GPSIMD microcode
                       │  ioctl(/dev/neuronN, ...)
                       ▼
  KERNEL       aws-neuronx-dkms  ── ioctl, DMA rings, DHAL v2/v3/v4, NQ, reset, pod
                       │  PCIe BAR MMIO · DMA · MSI-X · FW-IO mailbox
                       ▼
  ON-DEVICE    NeuronCore engines (PE/ACT/POOL/DVE/SP) + NCFW sequencer (Xtensa)
               + GPSIMD/Q7 vector cores (Vision-Q7)

How the data flows (the spine of this book)

A model's journey is the reading order: silicon model (Part I) → kernel driver (Part III) → runtime core (Part IV) → NEFF parse, memory plan, resource build (Part V) → TPB instruction lowering and ISA validation (Part VI) → execute: submit and harvest (Part VII) → DMA descriptors (Part VIII) → on-device collectives (Part IX) → collective firmware (Part X) and GPSIMD microcode (Part XI) → multi-node collectives (Part XII) → trace and telemetry (Part XIII). The cross-cutting apparatus — binary forensics (Part II), the datastore (Part XIV), security (Part XV), and reference tables (Part XVI) — wraps the spine. New readers should start with An Inference, End to End.

Evidence basis

The kernel driver ships as unstripped GPL source, and every binary in scope retains at least a full symbol table — several also carry DWARF — so the great majority of claims here are symbol- (and where DWARF survives, type-) grounded rather than pattern-matched:

• The DKMS C source is unstripped GPL-2.0 — kernel pages cite file:line directly.
• libnrt.so, libnds.a, and libnccom.so preserve DWARF — function names, struct field names, and enum values survive; pages cite addresses, struct offsets, and DWARF-recovered types. The two firmware carriers libncfw.so and libnrtucode_extisa.so (and libnccom-net.so) are unstripped but symtab-only — no DWARF; pages for those binaries cite named symbols and .rodata offsets, not DWARF types.
• The on-device firmware (NCFW Xtensa sequencer; GPSIMD Vision-Q7 microcode) is fully disassemblable: the Tensilica .tie config ships in the GPSIMD toolchain, so the custom vector ISA decodes exactly. See Part XI.

CORRECTION — an earlier scaffold of this index claimed all four runtime-lib binaries plus both collectives binaries "preserve DWARF". That is overturned by readelf -SW <bin> | grep -c '\.debug': only libnrt.so (9), libnds.a (835), and libnccom.so (8) carry .debug_* sections. libncfw.so (0), libnrtucode_extisa.so (0), and libnccom-net.so (0) are unstripped but symtab-only — no DWARF; this matches what all four firmware pages already state for libncfw.so. Pages for the symtab-only binaries are grounded on named symbols and .rodata offsets, not DWARF-recovered types.

CORRECTION — an earlier scaffold of this wiki described the GPSIMD/Q7 cores as "ARM-derived" with a TIE config that could not be decoded. Both are wrong: the GPSIMD compute cores are Tensilica Vision-Q7 (Xtensa LX with the IVP vector extension), and the .tie is shipped, so the 1065-op vector ISA is fully recovered (GPSIMD ISA Catalog). The separate FW-IO management path is documented at FW-IO MiscRAM Mailbox.

Companion wikis

• neuronx-cc — the compiler that produces the NEFFs nrt_load consumes
• neuronx-gpsimd — the GPSIMD/Q7 custom-op toolchain (the producer side of Part XI)
• neuron-jax-stack — the PJRT plugin that drives nrt_*
• neuronx-distributed — the distributed-training layer above the collectives
• neuronx-misc — diagnostic tools (neuron-monitor, neuron-ls, neuron-profile)