Userspace Runtime Core — Internal Architecture Map

All addresses, offsets, and symbol names on this page apply to libnrt.so from aws-neuronx-runtime-lib 2.31.24.0-0b044f4ce (real file libnrt.so.2.31.24.0, SONAME libnrt.so.1, build-id 8bb57aba0fb2e0035f1d88e9fc4fb3e7387c102e, git 0b044f4ce). The ELF is 122,956,336 bytes, ELF64 x86-64, not stripped, and carries full DWARF (.debug_info present; DW_AT_comp_dir = /opt/workspace/KaenaRuntime/build/private/develop/almalinux/RelWithDebInfo, DW_AT_producer = GNU C++17 14.2.1 -O2 -fPIC). All four PT_LOAD segments are identity-mapped, so .text, .rodata, and .data are VMA==fileoffset (read .data globals at their VMA directly); .bss is NOBITS. Other versions will differ.

CORRECTION — an earlier revision of this line claimed .data/.bss differ from their VMA by a 0x400000 delta. That is wrong for libnrt.so. readelf -lW shows the RW LOAD segment as LOAD 0xbeeaa0 0x0000000000beeaa0 0x0000000000beeaa0 0x049180 0x0c5860 RW 0x1000 (Offset == VirtAddr == 0xbeeaa0, delta zero), and readelf -SW shows .data at Address 0xc07e00 / Off c07e00 — identical. The 0x400000 delta is a fact about a different binary (the libtpu / Kaena-profiler image), not libnrt. Evidence grade: Confirmed (byte-anchored) — the layer-to-TU map is read from the binary's own 331-CU DWARF compile-unit table (P1-F-SRCMAP); every layer pins to a verbatim DW_AT_name source-TU path and at least one nm-verified symbol. · Part IV — Userspace Runtime Core · back to index

Abstract

libnrt.so is the userspace Neuron Runtime: the single shared object an ML framework (PyTorch-Neuron, JAX, the aws-neuronx-collectives stack) links against to load a compiled model (a NEFF) and run inference on Trainium/Inferentia silicon. It is one source tree — /opt/workspace/KaenaRuntime/ — compiled to 4,407 first-party functions across 203 translation units, with another 4,384 functions of statically-linked vendored code (Abseil, the Rust neuron_rustime runtime, the KaenaProfilerFormat protobuf, Libarchive+zlib, Simdjson, and the KaenaDriverLib shim). DWARF exposes the whole thing: there are no opaque blobs in the call path, only inlined headers. This page is the map of that tree — it names each architectural layer, pins it to the owning TU(s) and a representative symbol, shows how a runtime call descends the stack, and routes the reader to the ~26 sibling pages that derive each layer byte-for-byte.

The familiar reference frame is a classic layered userspace device runtime, the shape of a CUDA userspace driver or a vendor OpenCL ICD: a stable public C ABI at the top, a thin kernel-boundary syscall shim at the bottom, and — in between — an execution manager that owns the model database and the work queue, a device-driver core that owns device memory and the instruction-block builders, a vendored hardware-abstraction layer for register access, and a per-silicon device layer that specializes geometry and ISA for each chip generation. The runtime never issues a raw ioctl from anywhere except the bottom layer; every device transaction funnels through the ndl_* boundary into /dev/neuron<N>. That single funnel is the invariant a reimplementer must preserve.

The page is organized top-down. After the reimplementation contract and the at-a-glance table, §1 gives the six-layer stack diagram with the owning TU and symbol prefix for each layer; §2 walks one representative call (nrt_execute) down all six layers to show where each boundary is crossed; §3 is the layer → TU → owning-sub-page routing table that orients the rest of Part IV; and §4 is the cross-cutting infrastructure sidebar (logging, the device book, interned strings, virtual rings) that every layer leans on. Each layer's internal algorithm, struct layout, and decision logic live on a sibling page — this page links them and does not duplicate their derivations.

For reimplementation, the contract of the runtime core is:

• The six-layer descent — nrt_* public API → KMGR exec manager → TDRV device-driver core → KaenaHal register/platform adapter → NDL kernel-boundary shim → kernel ioctl. Each layer calls only the layer directly below it (plus the cross-cutting infra of §4); a reimplementation that lets, e.g., TDRV ioctl the kernel directly breaks the compatibility-mode fallback the NDL layer implements.
• The single kernel boundary — every device syscall (open/ioctl/mmap on /dev/neuron<N>, magic byte 'N') is emitted exclusively from ndl.c; the layers above hold only opaque handles (H_NN, dmem_t*, mem-handles) and never a raw fd except the one cached inside a mem-handle.
• The per-arch fan-out — three silicon generations (codenames sunda, cayman, mariana) each get their own tdrv_arch_*.c, instruction_block_*.c, and dve_dynamic_config_*.c; a one-shot arch-ops vtable (tdrv_arch_ops @0xc97180) selects the family at bring-up.
• The TU ownership map — which source TU implements which layer, so a reimplementer can mirror the module boundaries rather than reinvent them.

1. The Six-Layer Stack

The runtime core is six layers deep from the framework-visible ABI to the kernel syscall. Each layer is owned by a distinct family of TUs under /opt/workspace/KaenaRuntime/ (the bottom two are statically-linked vendored TUs). The diagram below pins each layer to its owning TU(s) and the symbol prefix that names its functions; the per-layer derivation lives on the sibling page named in §3.

 framework (PyTorch-Neuron / JAX / aws-neuronx-collectives)
        │  links against the @@NRT_2.0.0 versioned ABI
        ▼
 ┌─────────────────────────────────────────────────────────────────────────────┐
 │ [1] PUBLIC C API          nrt_*        nrt/*.cpp  (27 TU / 377 fn)            │
 │     init/load/exec/tensor/profile      nrt_init@0x94e90  nrt_execute@0x91de0  │
 │     • per-thread nlog tag • 4-state NRT_INIT_STATE guard • delegate, flush    │
 └─────────────────────────────────────────────────────────────────────────────┘
        │  nrt_execute → kmgr_*   ·   nrt_load → kmgr_load_nn_nc
        ▼
 ┌─────────────────────────────────────────────────────────────────────────────┐
 │ [2] EXEC MANAGER (KMGR)   kmgr_/dlr_   kmgr/dlr.cpp + kmgr/xu/* (10 TU/117 fn)│
 │     model DB (dlr_model_set) • sync XU work-queue • async worker engine       │
 │     kmgr_init@0xde080  kmgr_load_nn_nc@0xde280  dlr_kickoff_exec               │
 └─────────────────────────────────────────────────────────────────────────────┘
        │  schedule → build descriptors → ring doorbell
        ▼
 ┌─────────────────────────────────────────────────────────────────────────────┐
 │ [3] DEVICE-DRIVER CORE    tdrv_/dmem_/ tdrv/*.c  (84 TU / 2413 fn)            │
 │     hw_exec_queue/tensor   device bring-up • dmem allocator • instruction     │
 │     tdrv_init@0x26a310     blocks • DMA rings • per-arch arch_ops vtable       │
 │     arch_ops@0xc97180      tensor objects • hw_exec_queue • sync-point         │
 └─────────────────────────────────────────────────────────────────────────────┘
        │  al_reg_read32/write32 · al_mem_read/write_buf · al_copy_descriptor
        ▼
 ┌─────────────────────────────────────────────────────────────────────────────┐
 │ [4] KaenaHal ADAPTER      al_*/aws_hal_ tdrv/hal_platform.c + helper.c        │
 │     register/mem/barrier   al_reg_read32@0x2658a0 → csr.c (csr_read@0x315680) │
 │     ── backs vendored KaenaHal-2.31.0.0 (al_udma_/al_iofic_/aws_hal_*)        │
 └─────────────────────────────────────────────────────────────────────────────┘
        │  ndl_memory_alloc · ndl_memory_map · ndl_memory_copy_as · ndl_*_ioctl
        ▼
 ┌─────────────────────────────────────────────────────────────────────────────┐
 │ [5] NDL kernel shim       ndl_*        KaenaDriverLib-2.27.4.0/src/ndl.c (123)│
 │     IOCTL/mmap wrappers    ndl_device_ioctl (per-dev fd) · ndl_misc_ioctl     │
 │     + compatibility-mode (_cm) fallback stubs for missing driver features     │
 └─────────────────────────────────────────────────────────────────────────────┘
        │  open()/ioctl(magic 'N')/mmap() on /dev/neuron<N>
        ▼
 ╔═════════════════════════════════════════════════════════════════════════════╗
 ║ [K] KERNEL DRIVER (DKMS, out of this binary) — Part III · ncdev_ioctl dispatch║
 ╚═════════════════════════════════════════════════════════════════════════════╝

 cross-cutting (called from every layer — §4):
   nlog/  (logging+backtrace)   ·   tdrv/db.c+dbtc.c  (device book)
   nrt/nrt_interned_string_db.cpp (interned strings)  ·  tdrv/vring.c (virtual rings)

NOTE — layers [4] and [5] are vendored TUs statically linked in, not first-party KaenaRuntime code, but the boundary they form is real and load-anchored. NDL's source package (KaenaDriverLib-2.27.4.0) lags the runtime package (2.31.24.0); this version skew is a confirmed DWARF fact, not an extraction error. The KaenaHal package (KaenaHal-2.31.0.0) has no DWARF CU of its own in libnrt.so — its al_udma_*/aws_hal_* code reaches the binary only through inlined headers; the al_* adapter that backs it (tdrv/hal_platform.c) is first-party and is what layer [4] documents.

Why the layering is shaped this way

The split is not cosmetic; each boundary earns its place.

• API ↔ KMGR separates the idempotent, fork-aware, state-guarded public surface (every nrt_* entry runs the same NRT_INIT_STATE preamble and a tail log-flush — see api-lifecycle) from the model-database orchestration that does not care about C-ABI concerns. KMGR keys everything on dlr_model_set (an unordered_map<u32 H_NN.id, dlr_model*> under an rwlock), so the API layer can hand out opaque H_NN handles and never touch model internals.
• KMGR ↔ TDRV separates which model to run and when from how the device is driven. KMGR builds and rings; TDRV owns every device resource (the dmem allocator, the per-arch instruction-block builders that are the single densest code mass at 783 combined functions, the DMA rings, the tensor objects).
• TDRV ↔ KaenaHal is the Alpine porting seam: TDRV calls the Annapurna/Alpine al_* register and memory primitives; the adapter (hal_platform.c) implements those primitives over csr.c (PCIe CSR access) so the vendored al_udma/al_iofic/aws_hal code compiles unchanged.
• KaenaHal ↔ NDL ↔ kernel is the syscall funnel. NDL is the only TU that opens /dev/neuron<N> and issues ioctl; it also implements the compatibility-mode fallback (the _cm stub set) so that a runtime built against a newer ABI degrades gracefully on an older driver. Letting any higher layer syscall directly would defeat that fallback.

2. How a Call Descends the Stack

The clearest way to read the stack is to follow one call. Below is nrt_execute — the hot inference path — from the framework's call to the kernel doorbell. Each → crosses one layer boundary; the symbol that does the crossing is named. (The full per-stage logic is owned by exec/kmgr-facade and exec/submit-path; this is the descent skeleton only.)

// [1] PUBLIC API — nrt/nrt_exec.cpp
nrt_execute(model, in_tensors, out_tensors):     // @0x91de0
    nlog_coalescing_init_thread()                //  per-thread log tag
    if NRT_INIT_STATE != INIT: return guard_status   //  4-state guard (api-lifecycle.md)
    model_id = nrt_get_interned_model_id(...)    //  @0x96fa0 → interned-string DB (§4)
    rc = kmgr_sync_exec(model->h_nn, io, ...)    //  → [2] cross into exec manager
    nlog_flush_if_unflushed("Generic API Failure"); return rc

// [2] EXEC MANAGER — kmgr/dlr.cpp + kmgr/xu/*
kmgr_sync_exec(h_nn, io):                        //  resolve dlr_model from dlr_model_set (rwlock + refcount)
    req = tpb_xu_schedule_request(model, io)     //  kmgr/xu/tpb_xu.cc — build a work-queue request
    dlr_add_to_hw_exec_queue(req)                //  → [3] the shared "build + ring" funnel
    dlr_kickoff_exec(req)                         //  → [3]
    kmgr_exec_wait(req)                           //  eventfd-block on completion (NQ harvest)

// [3] DEVICE-DRIVER CORE — tdrv/hw_exec_queue.c + tdrv/dma_memory.c + tdrv/tensor.c
hw_exec_queue_add_exec_request(req):             //  bind tensors, stage inputs
    dmem_buf_batch_copyin(io.inputs)             //  tdrv/dma_memory.c @0x228eb0 — H2D staging
        → dmem_batch_transfer_common(...)        //  build neuron_memcpy_batch_t[]
            → ndl_memory_buf_batch_copy(...)     //  → [5]  (one batched copy ioctl)
    instruction_block_<arch>_build(...)          //  per-arch IB emitter (sunda/cayman/mariana)
    csr_write(doorbell_reg, ring_tail)           //  → [4]  ring the hardware doorbell

// [4] KaenaHal ADAPTER — tdrv/hal_platform.c → tdrv/csr.c
al_reg_write32(reg, val):                        //  @ hal_platform.c — porting-shim primitive
    csr_write(bar, off, val)                     //  tdrv/csr.c @0x315820 — PCIe CSR write
        → (mmap'd BAR store; mapping established earlier via ndl_memory_map → [5])

// [5] NDL KERNEL SHIM — KaenaDriverLib-2.27.4.0/src/ndl.c
ndl_memory_buf_batch_copy(device, batches, n, dir):
    fd = device->context.nd_fd                   //  cached per-device fd (struct +632)
    ioctl(fd, NEURON_IOC_<copy>, &arg)           //  → [K]  magic 'N' (0x4E); -errno on fail
                                                 //  or, if driver lacks the ioctl: *_cm fallback stub

Three facts a reimplementer must lift from this descent:

• The kernel boundary is reached only at [5]. The H2D copy in [3] does not ioctl; it builds a neuron_memcpy_batch_t[] array and hands it to one ndl_memory_buf_batch_copy — a single batched syscall, not one per tensor. The doorbell ring in [4] is a store to an already-mmap'd BAR, and that mapping was itself established earlier by ndl_memory_map at [5].
• Layer [4] is a porting shim, not a dispatcher. al_reg_write32 adds no arch logic; it forwards to csr_write. The arch-specific behavior lives one layer up, in the tdrv_arch_ops vtable and the per-arch instruction_block_* emitters at [3].
• The compatibility-mode fork is at [5], invisible above. Whether the batched copy uses the real ioctl or the _cm fallback is decided inside ndl.c based on driver capability; layers [1]–[4] are oblivious. (See ndl for the dual-dispatch pattern.)

3. Layer → TU → Owning Sub-Page

This is the routing table for the rest of Part IV: each architectural layer, the TU family that implements it (verbatim /opt/workspace/KaenaRuntime/ paths from the DWARF CU set), a representative nm-verified symbol, and the sibling page that derives it. Confidence is HIGH throughout — every TU path is a verbatim DW_AT_name string and every symbol is nm-verified (P1-F-SRCMAP §6).

QUIRK — the public API is the smallest layer by code mass (377 fns across 27 TUs) and the device-driver core is by far the largest (2,413 fns across 84 TUs). Most nrt_* public functions are thin wrappers — nrt_tensor.cpp (5 fns) forwards straight to tdrv/tensor.c (26 fns), nrt_async_sendrecv.cpp thunks into enc/async_sr/*. A reimplementer who budgets effort by the public-header surface will badly under-scope TDRV. The weight is in the instruction-block builders and the collectives engine (Part IX), not the ABI.

NOTE — the on-device collectives engine (enc/ + encd/, 52 TUs / 1,062 fns) and the NEFF/KELF container loader (kelf/, 5 TUs / 318 fns, with Libarchive+zlib+Simdjson) are peer subsystems that sit beside this six-layer spine, not inside it — they are driven from KMGR/TDRV but own their own architecture. They are mapped by Part IX and Part V respectively and are out of scope for this page.

3.1 The per-arch fan-out

Three of the layers above are triplicated by silicon generation, and the triplication is a structural fact a reimplementer must mirror, not flatten. The DWARF CU set shows the same three codenames recurring across TDRV, the per-arch device layer, and the DVE config: sunda, cayman, mariana. Each gets its own copy of the geometry/ISA/config TUs, and a single one-shot vtable picks the family at bring-up.

The selector is the tdrv_arch_ops vtable (@0xc97180, .bss, 488 B / 61 physical 8-byte slots (47 named members)), populated exactly once, lazily, by one of tdrv_arch_register_{sunda,cayman,mariana} according to the silicon's arch type (gate tdrv_arch_ops_initialized @0xc97368; switch(arch_type){2→sunda, 3→cayman, 4→mariana}). Every tdrv_arch_get_* geometry accessor double-checks the gate and routes through a slot — so the per-arch divergence is resolved by one indirect call, not by if (arch == …) scattered through the device-driver core. The instruction-block emitters are the exception: they are dispatched directly by the IB builder rather than through the vtable, and they are the single densest code mass in the binary (783 combined functions). The arch-ops install and the geometry accessors are owned by tdrv-arch-ops; the per-arch instruction builders by isa/validators-per-arch.

CORRECTION — the slot count is standardized on tdrv-arch-ops: the table is 488 B / 61 physical 8-byte slots (47 named members) (gap 0xc97368-0xc97180 = 0x1e8 = 488 = 61 slots). The "47 slots" phrasing undercounts the physical layout by 14 — the 47 figure is the count of named members, not slots.

GOTCHA — the codenames are silicon generations, not products: a reimplementer mapping them to cloud instance names must go through the arch enum (arch_type 2/3/4), not the marketing name. The mapping is owned by arch/generations-enum. Do not assume a fourth …_v4 arch exists in this build merely because the enum reaches 4 — instruction_block_mariana.c inlines its isa.c, so the per-arch TU inventory is asymmetric (sunda/cayman have a separate isa.c, mariana does not).

4. Cross-Cutting Infrastructure

Four small subsystems are called from every layer of the stack and belong to no single one. They are the runtime's shared plumbing; the stack diagram of §1 lists them at the bottom as cross-cutting. Each has its own derivation page; this sidebar names the TU, the entry symbol, and the role so the reader knows where each utility lives.

Two of these are worth a reimplementer's attention up front. The device book is the indirection that lets every layer name a NeuronCore without holding device pointers: dmem_alloc_internal calls db_physical_core_get_mla_and_tpb(pcore, &mla, &tpb) to resolve the owning MLA/TPB before it allocates, and db_tdrv_ctx_get() returns the per-process tdrv_ctx_t that the profile and EFA-BDF paths read. The interned-string DB is what makes the trace plane cheap: a model id is hashed to a uint64 once (nrt_get_interned_model_id @0x96fa0 → nrt_interned_string_db_get_id) and thereafter every trace event carries the integer, not the string. Both are global singletons initialized inside nrt_init's subsystem-init phase (nlog_init, then the per-vcore book) and are torn down in reverse by nrt_close.

NOTE — the trace/profile producer chain (nrt/nrt_profile.cpp, nrt_sys_trace*.cpp, nrt_inspect*.cpp) and its vendored KaenaProfilerFormat protobuf backend (ntff.pb.cc/neuron_trace.pb.cc, 730 generated fns) read these infra services heavily but are their own subsystem — see Part XIII. The Rust neuron_rustime capture engine bridges into nlog and the interned-string DB via a cbindgen FFI boundary (trace/rust-ffi).

Cross-References

The six layers, top to bottom

• api-lifecycle — layer [1]: nrt_init/nrt_close, the 4-state NRT_INIT_STATE guard, driver-compat-id negotiation
• api-device-config — layer [1]: device/topology query wrappers and config parsing
• exec/kmgr-facade — layer [2]: the KMGR exec manager facade and the dlr_model_set model DB
• tdrv-lifecycle — layer [3]: tdrv_init device bring-up and the tdrv_arch_ops vtable
• tdrv-dmem — layer [3]: the dmem device-memory allocator (the worked-example H2D staging path)
• hal-adapter — layer [4]: the al_* KaenaHal platform-services adapter over csr.c
• ndl — layer [5]: the ndl_* IOCTL/mmap shim and the compatibility-mode fallback dispatch

Cross-cutting and peer subsystems

• device-book · interned-strings · logging — the §4 shared plumbing
• tdrv-tensor — the nrt_tensor_t object layer that bridges API [1] and TDRV [3]
• collectives/overview — the enc/encd peer subsystem beside the spine (Part IX)
• neff/overview — the NEFF/KELF container loader peer subsystem (Part V)
• kernel/ioctl-dispatch — the kernel ncdev_ioctl side of the NDL boundary (Part III)
• front/source-tree — the full 203-TU KaenaRuntime source-tree reconstruction this map summarizes
• back to index