Per-Arch Instruction Validators (SUNDA / CAYMAN / MARIANA Deltas)

All addresses, offsets, struct ordinals, and enum values on this page apply to libnrt.so from aws-neuronx-runtime-lib_2.31.24.0-0b044f4ce (runtime-lib 2.31.24.0-0b044f4ce; libnrt.so.2.31.24.0, SONAME libnrt.so.1, ELF64, not stripped, DWARF present, build-id 8bb57aba0fb2e0035f1d88e9fc4fb3e7387c102e). .text VMA equals file offset for the cited ranges; the per-arch validator bands are all in .text, so every band/validator address below is a direct file offset. Source-TU attribution is from DWARF addr2line; function addresses/sizes/caller-counts from the function table (functions.json); the dispatcher jump tables from switches.json; struct/enum layouts from DWARF (structures.json, enums.json). Other versions will differ.

Abstract

The validator architecture page documents one validator forest — the silent is_valid_<op> gate, its dbg_is_valid_<op> twin, the shared has_valid_neuron_header/has_valid_neuron_events preamble, and the NEURON_ISA_TPB_DEBUG_BOOL return ABI. This page documents the fact that that entire forest exists three times over. libnrt.so carries a complete, independent copy of the validator library for each of the three NeuronCore silicon generations it supports — SUNDA (arch 2, ISA-v2), CAYMAN (arch 3, ISA-v3), and MARIANA (arch 4, ISA-v4) — each occupying its own contiguous .text band of hundreds of kilobytes, each emitted from its own per-arch copy of the ISA-spec header. The three copies are not a runtime polymorphism over one body; they are three separately-compiled bodies, selected statically per record by the arch-versioned entry ib_is_valid_engine_instruction_v{2,3,4} (the tdrv_arch_ops slot final_inst_validity_check, runtime/tdrv-arch-ops §4.3). This is the single most expensive structural decision in the validation layer, and the one a reimplementer most needs to understand before deciding whether to replicate it.

The triplication is generated, not hand-maintained. DWARF addr2line attributes every validator in the SUNDA band to /opt/amazon/include/neuron_sunda_arch_isa/tpb/aws_neuron_isa_tpb_assert.h, every validator in the CAYMAN band to the identical path under neuron_cayman_arch_isa/, and every validator in the MARIANA band to neuron_mariana_arch_isa/. The three headers are the same generated file parameterized by one arch token: for any opcode, the SUNDA is_valid_<op>, the CAYMAN is_valid_<op>_0, and the MARIANA is_valid_<op>_1 are the same predicate tree compiled against three different sets of legality constants (address apertures, quadrant geometry, active-channel windows, the dtype set, and the small handful of per-generation opcodes). The IDA _0/_1 suffixes and .constprop.0/.part.0 decorations are GCC clone/constprop artifacts, not part of the C identifier; nm shows the clean symbol. The cost is paid in .text: the same dbg_is_valid_dtype body is compiled three times, the same tensor4d_valid is compiled three times, and the giant silent dispatcher exists as is_valid_neuron_instruction @0x2f2350 (17 952 B, SUNDA), is_valid_neuron_isa_instruction @0x3a60b0 (16 606 B, CAYMAN), and a third copy inlined into ib_is_valid_engine_instruction_v4 @0x43f810 (16 087 B, MARIANA).

This page treats the triplicated space by its two axes — opcode-class × arch — and never as three per-arch dumps. The roster (§2) lists each validator family once, with an arch-coverage column, because the families are shared; the delta table (§3) is the centerpiece: opcode-class rows against {SUNDA v2, CAYMAN v3, MARIANA v4} columns, with each cell recording present / absent / field-delta rather than re-listing the (identical) check tree per arch. A reimplementer who internalizes the two axes can reconstruct any one arch's validator set from the shared family roster plus that arch's column, instead of memorizing three near-identical rosters. §1 fixes the generation scheme and pins the three band extents; §4 documents the ib_hal_to_isa_engine_v{3,4} HAL→ISA engine map that sits at the head of each arch's block driver and routes a record to its per-engine legality switch.

For reimplementation, the contract is:

• The triplication invariant — there is one validator design and three compiled instances. For every opcode, {SUNDA, CAYMAN, MARIANA} each carry an independent is_valid_<op> + dbg_is_valid_<op> pair in their own band, generated from neuron_<arch>_arch_isa/tpb/aws_neuron_isa_tpb_assert.h. Reproduce the design once; instantiate it per target with that target's constant table.
• The static arch selection — the arch is chosen at the entry, not inside the body. tdrv_arch_ops.instr_block_caps->final_inst_validity_check points at ib_is_valid_engine_instruction_v2/_v3/_v4, latched once per process by al_hal_tpb_get_arch_type() ∈ {2,3,4}. Each entry calls only its own arch's forest; the three forests never cross-call.
• The shared family roster — the set of validator families (header/events preamble, the tensorNd_valid descriptor ladder, tensor_start_addr_valid/addr_aligned_dtype, is_valid_dtype/is_valid_enum, the per-opcode validators) is the same across all three arches; it is the rows of §2 and §3.
• The per-arch deltas are constants and a few opcodes, not logic — the arches differ in address-aperture and quadrant geometry constants, the dtype set (FP4 / 64-bit gating), and a small per-generation opcode delta (e.g. CAYMAN gains dma_transpose 0xBD, jpeg_decode 0x81; MARIANA gains the dma_gather_transpose, modify_pool_config, FP4 quantize paths). The opcode→struct dispatch and the predicate trees are identical.

1. The Generation Scheme — One Template, Three Compiled Copies

Purpose

Before reading any delta, a reimplementer needs the mechanical fact that makes the deltas small: the three validator forests are not three hand-written validators that happen to resemble each other — they are three compilations of one generated source. The source is a per-arch copy of aws_neuron_isa_tpb_assert.h; the compiler emits, for each opcode, a silent predicate and a debug twin into each arch's translation unit, and the linker leaves them as three disjoint .text bands. That is why the per-opcode check trees are structurally identical across arches and the only differences are constant values and a few generation-specific opcodes.

The template → three TUs mechanism

                          aws_neuron_isa_tpb_assert.h   (one generated template)
                                          │
              ┌───────────────────────────┼───────────────────────────┐
   neuron_sunda_arch_isa/tpb/   neuron_cayman_arch_isa/tpb/   neuron_mariana_arch_isa/tpb/
        (arch token = sunda)        (arch token = cayman)        (arch token = mariana)
              │                           │                           │
        compiled TU                 compiled TU                 compiled TU
              │                           │                           │
   ┌──────────┴──────────┐     ┌──────────┴──────────┐     ┌──────────┴──────────┐
   SUNDA .text band          CAYMAN .text bands          MARIANA .text bands
   is_valid_<op>             is_valid_<op>_0 (clone)     is_valid_<op>_1 (clone)
   dbg_is_valid_<op>         dbg_is_valid_<op>_0         dbg_is_valid_<op>_1
   tensorNd_valid            tensorNd_valid_0            tensorNd_valid_1
   (~0x27xxxx, 0x33xxxx)     (0x387540.., 0x39c610..)    (0x3b0960.., 0x3e0980..)

The _0 / _1 symbol suffixes on the CAYMAN and MARIANA copies are GCC clone disambiguators (the SUNDA copy, emitted first, keeps the undecorated name); .constprop.0 and .part.0 are constant-propagation and partial-inlining decorations. nm resolves all of them to the clean C identifier. The structural proof that the three are one design is DWARF: each band's functions addr2line to the same relative path tpb/aws_neuron_isa_tpb_assert.h, differing only in the neuron_<arch>_arch_isa/ prefix, and the line numbers within that header match across arches (e.g. tensor4d_valid resolves to assert.h:3089 in the MARIANA copy at 0x3e6af0).

Band extents (pinned)

The three forests do not interleave — each is a run of contiguous .text. The documented extents below are the slices verified function-by-function against nm (band boundaries confirmed by the first OUT-of-band symbol):

NOTE — the bands are adjacent and partly intermixed at the seams, not perfectly disjoint. The tail of the CAYMAN band 0x3aad80..0x3ab42e holds five MARIANA shared primitives (neuron_isa_tpb_new/merge_dbg_bool_1, is_valid_enum_1, is_valid_dtype_1, is_valid_int_aluop_0) physically abutting the CAYMAN block — they are high-degree boundary nodes fanned-in by 49–107 MARIANA _1 validators living in other bands. DWARF attribution, not address proximity, is the ground truth for which arch a function belongs to. A reimplementer sweeping by address range must use addr2line to assign each function to its arch, or it will mis-file these seam functions.

QUIRK — the three giant silent dispatchers are the clearest single artifact of the triplication cost. is_valid_neuron_instruction @0x2f2350 (SUNDA, 17 952 B, 530 BB), is_valid_neuron_isa_instruction @0x3a60b0 (CAYMAN, 16 606 B, 499 BB), and the third copy folded into ib_is_valid_engine_instruction_v4 @0x43f810 (MARIANA, 16 087 B, 488 BB) decode the same 64-byte record with the same opcode→struct switch — the opcode dispatch is byte-for-byte identical across all three. Only the leaf legality constants the arms feed differ. Roughly 50 KB of .text implements one logical function three times.

2. The Validator-Family Roster

Purpose

The triplicated space is best described by listing each validator family once and recording which arches carry it. Because the families are shared (they come from the one template), this single roster — not three per-arch lists — is the complete inventory. The arch-coverage column flags the handful of families that are generation-specific; everything unmarked exists, identically-structured, in all three bands.

Roster (one row per family, all three arches)

GOTCHA — the per-opcode validator names drift slightly between the silent and debug twin and between arch clones (is_valid_s2s2d3_roi vs dbg_is_valid_s2s2d3_roi_0; tensor2d_valid vs tensor2d_valid.constprop.0 vs tensor2d_valid_1). These are the same family — do not count them as distinct validators. A reimplementer who treats the silent leaf, its debug twin, and the three arch clones as five separate functions will over-count the design five-fold; there is one predicate per (opcode-class), instantiated 2 (silent/debug) × 3 (arch) ways.

3. The Per-Arch Delta Dimension Table

This is the centerpiece. The validator space has two axes: opcode-class (the rows) and arch (SUNDA v2 / CAYMAN v3 / MARIANA v4, the columns). Every cell records the delta — present (the family exists, structurally identical), +field-Δ (present but with an arch-specific constant or extra field check), or absent (the opcode/check does not exist on that generation). The table deliberately does not repeat the (identical) check tree per column; it records only how each opcode-class differs across the three silicon targets. A representative validator address pins each present cell to its band so the claim is verifiable.

Axis values

Dimension table (opcode-class × arch)

QUIRK — read the table by column delta against SUNDA, not by row. The SUNDA column is the baseline; CAYMAN's column adds the copy/cast/reciprocal, jpeg/sort/conv-lut, and dma_transpose opcode-classes (the 0x46-band engine ops and the DMA-transpose engine), and MARIANA's column further adds the FP4 dtype gate, dma_gather_transpose, dma_indirect 0xBB, the modify_pool_config/ctrl_es/ctrl_poll_sem composites, and the s3s3d1_bn/bn2 batchnorm opcodes. Every unmarked cell is the same predicate tree — the deltas are a small additive frontier of new opcodes per generation plus a handful of +Δ const aperture/quadrant geometry changes. A reimplementer reproduces the SUNDA forest, then adds CAYMAN's new opcode validators, then MARIANA's; the shared core is never re-derived.

NOTE — the +Δ const cells (address aperture, quadrant geometry, active-channel windows) are where the per-silicon physical SBUF/PSUM layout leaks into the validator. The check logic is identical — (addr & mask) within aperture — but the mask and aperture constants are per-generation, because the partition stride, PSUM bias, and quadrant size are hardware geometry. These constants are the validator's encoding of the device address map; they are not free to copy between arches.

4. The ib_hal_to_isa_engine HAL→ISA Engine Map

Purpose

Each arch's block driver opens by mapping the runtime's HAL engine enumeration onto the ISA's NEURON_ISA_TPB_NEURON_ENGINE taxonomy, then runs a per-engine opcode-legality switch (a PE-array matmul issued on the SP engine is malformed regardless of field validity). The map is a tiny per-arch leaf — ib_hal_to_isa_engine_v3 @0x3aa190 (72 B, instruction_block_cayman.c:33) for CAYMAN, and its MARIANA analog _v4 @0x43f7c0 reached from ib_is_valid_engine_instruction_v4 @0x43f810 (instruction_block_mariana.c:33). It is the one place the otherwise-generated validator forest meets hand-written tdrv glue, and it is the entry seam at which arch selection has already happened.

Entry Point

tdrv_arch_ops.instr_block_caps->final_inst_validity_check   ── slot, arch-latched
  └─ ib_final_inst_validity_check_v3   (0x3aaca0)  ── 64-byte-stride block driver
       └─ ib_is_valid_engine_instruction_v3 (0x3aa1e0)      ── per-record, engine-gated
            ├─ ib_hal_to_isa_engine_v3       (0x3aa190)     ── HAL eng → ISA eng  [THIS §]
            ├─ is_valid_neuron_isa_instruction (0x3a60b0)   ── structural dispatch (CAYMAN forest)
            └─ ib_is_valid_engine_instruction_v3_0 (0x3a4cf0) ── .part.0 logging tail
                 └─ debug_invalid_neuron_isa_instruction_0 (0x39fec0) ── DEBUG_BOOL reason

For MARIANA the same shape is ib_final_inst_validity_check_v4 @0x4436f0 → ib_is_valid_engine_instruction_v4 @0x43f810 → ib_hal_to_isa_engine_v4 @0x43f7c0 + the v4 forest, with the diagnostic tail in debug_invalid_neuron_engine_instruction @0x439ab0. The MARIANA driver is reached through the tdrv_arch_ops vtable slot 0xbf3778 (ib_final_inst_validity_check_v4_wrapper @0x30d2d0, tdrv_arch_mariana.c:62).

Algorithm

// Models ib_hal_to_isa_engine_v3 @0x3aa190 (72 B, instruction_block_cayman.c:33).
// The MARIANA copy ib_hal_to_isa_engine_v4 @0x43f7c0 is the same body, v4-cloned.
// Maps al_hal_tpb_eng_type_t (HAL taxonomy) -> NEURON_ISA_TPB_NEURON_ENGINE (ISA taxonomy).
function ib_hal_to_isa_engine_v3(hal_eng) -> NEURON_ISA_TPB_NEURON_ENGINE:
    if hal_eng > AL_HAL_TPB_ENG_SP:                  // out-of-range HAL id
        nlog_write("Unknown hal->isa engine enum conversion (%u --> ?)", hal_eng)  // 0x812f58
        return NEURON_ISA_TPB_TOP_SP                 // fallback: the collective-trigger sequencer
    return isa_engine_of[hal_eng]                    // direct table map; {PE,ACT,POOL,DVE,SP}->{0,1,2,3,4}

// Models ib_is_valid_engine_instruction_v3 @0x3aa1e0 (2747 B). The per-record gate that
// sits above the generated forest; one body per arch (v2/v3/v4), differing only in which
// arch-clone of the forest and which per-engine opcode set it consults.
function ib_is_valid_engine_instruction_v3(hal_eng, inst /* 64-byte record */) -> bool:
    isa_eng = ib_hal_to_isa_engine_v3(hal_eng)        // 0x3aa190 — HAL->ISA engine id
    if !is_valid_neuron_isa_instruction(inst):        // 0x3a60b0 — STRUCTURAL gate (CAYMAN forest)
        return ib_is_valid_engine_instruction_v3_0(inst)  // 0x3a4cf0 — log reason, return false
    // per-engine opcode legality: is this opcode issuable on THIS engine?
    switch isa_eng:                                   // {PE, ACT, POOL, DVE, SP}
        case PE:   ... is_valid_int_aluop / tensor_tensor_bitvec_engine / tensor_start_addr_valid_0
        case ACT:  ...                                // per-engine opcode-set predicates
        case POOL: ...                                // (bittest64 masks over the opcode byte)
        case DVE:  ...
        case SP:   ...
    return engine_opcode_legal                        // AND of structural + per-engine legality

// Models ib_final_inst_validity_check_v3 @0x3aaca0 (207 B). The block driver: walk the
// instruction-block buffer in 64-byte strides, validate each record, log the first failure.
function ib_final_inst_validity_check_v3(buf /* {size_used, *buf} */, hal_eng) -> bool:
    assert(buf.size_used % TPB_INST_NBYTES == 0)      // 0x40 stride invariant; 0x813000 assert
    for off in range(0, buf.size_used, 0x40):         // one 64-byte record per step
        inst = buf.buf + off
        if !ib_is_valid_engine_instruction_v3(hal_eng, inst):
            nlog_write("Found invalid instruction! opcode=%u, offset=%u",  // 0x813030
                       inst.raw.bytes[0], off)
            return false                              // first invalid record stops the walk
    return true

Function Map

Considerations

The engine map and the per-engine legality switch are the only hand-written tdrv code in the per-arch validation path; everything below is_valid_neuron_isa_instruction is generated. The nlog_write strings ("Unknown hal->isa engine enum conversion (%u --> ?)" @0x812f58, "Found invalid instruction! opcode=%u, offset=%u" @0x813030, "Error message number:%u, message: %s" @0x812f30) live only in these ib_* glue functions — the generated validators reference no strings, because they assemble their DEBUG_BOOL messages from the shared .rodata template pool (xmmword_850FA0.., validator-architecture §2). The provenance string /opt/workspace/KaenaRuntime/tdrv/instruction_block_cayman.c (@0x81f090) appears four times, once per ib_* call site in the CAYMAN driver — the cleanest confirmation that the driver layer is per-arch tdrv glue wrapping the per-arch generated forest.

QUIRK — the triplication is complete down to the block driver, not just the leaves. There is a separate ib_final_inst_validity_check_v{2,3,4}, a separate ib_is_valid_engine_instruction_v{2,3,4}, and a separate ib_hal_to_isa_engine_v{2,3,4} per arch — even these tiny glue functions (72–207 B) are cloned three times rather than shared with a runtime arch parameter. The arch is resolved once, statically, at the tdrv_arch_ops slot, and from that point every function on the path is the arch-specific clone. A reimplementer can collapse the three glue clones into one parameterized function with no behavioral change — the cloning here is a compilation artifact of the per-arch TU split, not a design requirement, since the glue logic (unlike the leaf constants) does not actually vary by arch.

Related Components

Cross-References

• ISA Validator Architecture and Entry Tree — the validator engine this page deltas: the silent/debug two-family design, the NEURON_ISA_TPB_DEBUG_BOOL ABI, and the shared header/events preamble (do not re-derive it here)
• The 64-Byte Instruction Record Format — the NEURON_ISA_TPB_INST_UNION and per-opcode struct views every per-arch validator decodes
• Overview: the TPB Engine Instruction Model — Part VI map; where per-arch validation sits in the lowering→validation→stage flow
• TDRV: arch-ops Dispatch and Sync-Event Accessors — the tdrv_arch_ops vtable (final_inst_validity_check slot, instr_block_caps) and al_hal_tpb_get_arch_type that statically select the arch forest