Cross-Generation Opcode-Table Diff + TONGA

This is the ISA-evolution step-diff for the GPSIMD opcode namespace and the definitive TONGA V1 deep-dive. It answers, opcode by opcode, what every generation step adds, changes, and removes across the unified ladder — SUNDA (v2) → CAYMAN (v3) → MARIANA (v4) → MARIANA_PLUS (v4+) → MAVERICK (v5) — with TONGA (V1) characterized as the pre-unified outlier from its own, separate ISA package. It carries four load-bearing results, each proven from the shipped bytes:

1. The opcode byte is a write-once global namespace. Every opcode that lives in two generations carries the identical hex value in both — proven across the full SUNDA↔MAVERICK span (138 common opcodes, zero value drift). New generations only append; the single re-baseline (v2→v3) abandons seven bytes that are never reused.
2. There are four ISA opcode tables, not five. NEURON_ISA_TPB_OPCODE exists in exactly the SUNDA / CAYMAN / MARIANA / MAVERICK trees. MARIANA_PLUS reuses the MARIANA table 1:1 — it is a firmware/register flag-refresh, not a distinct ISA.
3. TONGA V1 has a real opcode ISA of its own — an engine-scoped TONGA_ISA_TPB_OPCODE (byte = local | (ENGINE<<5)) with a 64-byte instruction-struct family — but not the unified flat namespace. V1 is the genesis of the v2+ opcode bytes (the datapath bytes are inherited verbatim) and the genesis of the flattening (v2 drops the engine bitfield).
4. There is no Vision-Q7 GPSIMD POOL engine in V1. V1's POOL is hardwired MaxPool/AvgPool; GPSIMD (the programmable NX_POOL Q7 cluster the rest of this book documents) is a SUNDA-onward (v2+) feature.

Every value, count, and presence cell below was read this pass from the shipped header enums (bracket-bounded), cross-checked against the per-gen instruction_mapping.json, and grounded with direct rg/bracket counts — never a decompile. Confidence/evidence tags follow The Confidence & Walls Model: HIGH/MED/LOW × OBSERVED/INFERRED/CARRIED.

This page owns the evolution narrative + TONGA. For the image-grounded, per-generation KIT→kernel/engine/struct Rosetta, see the sibling Cross-Gen Kernel-Info Matrix; for the per-header structural diff see Arch-ISA Header Diff; for the completeness boundary (140 real HW opcodes) see The Opcode Catalog Ledger; for the portability proof that justifies decoding one machine see The Gen-Invariance Thesis; for the (arch_id, coretype) identity algebra see The Codename → Generation Map.

1. What actually ships — the header-tree inventory

There are four NEURON_ISA_TPB_OPCODE enum trees (the authoritative opcode source), one TONGA V1 ISA package (a different opcode form), and zero mariana_plus opcode tree. [HIGH/OBSERVED]

All four counts are independently bracket-bounded (the count of NEURON_ISA_TPB_OPCODE_… = 0x… lines between typedef enum NEURON_ISA_TPB_OPCODE { and } NEURON_ISA_PACKED NEURON_ISA_TPB_OPCODE;), and each count includes the INVALID = 0xff sentinel. The path stem is …/custom_op/c10/include/neuron_<gen>_arch_isa/tpb/aws_neuron_isa_tpb_common.h. [HIGH/OBSERVED]

NOTE — MARIANA_PLUS has no opcode tree. There is no neuron_mariana_plus_arch_isa/ directory anywhere in the package. A register tree arch-headers/mariana_plus/ does exist, but it carries CSR/register-block headers (aws_mariana_isa_common.h, a2i, cce, address_map, …) with no aws_neuron_isa_tpb_common.h and no NEURON_ISA_TPB_OPCODE. MARIANA_PLUS executes the MARIANA 159-opcode table verbatim — see §7. [HIGH/OBSERVED]

1.1 The version axis is cumulative, and it is born at v2

NEURON_ISA_TPB_NEURON_CORE_VERSION grows with the tree, and no tree has a V4_PLUS / V4P enumerator: [HIGH/OBSERVED]

The SUNDA enumerator carries the decisive in-band comment:

typedef enum NEURON_ISA_TPB_NEURON_CORE_VERSION {
    NEURON_ISA_TPB_NEURON_CORE_VERSION_V2 = 2,   // V1 ISA is maintained in separate package
} ... ;

aws_neuron_isa_tpb_common.h:132 (sunda).

This single line is the structural keystone of the whole page: the version axis starts at v2, it has no V1 point (V1 is externalized by name to a separate package — that is TONGA, §5), and it has no v4+ point (MARIANA_PLUS is distinguished by a runtime string selector, not an ISA-version number — §7). [HIGH/OBSERVED]

2. The master opcode-presence matrix

Columns: V2=SUNDA, V3=CAYMAN, V4/4+=MARIANA (≡ MARIANA_PLUS, a single shared column — §7), V5=MAVERICK. Cells show the hex value where present, · where absent. TONGA (V1) is not a per-row column — its bytes are mappable onto this namespace (the V1 datapath bytes equal the v2 bytes, §6), but V1 uses a different encoding and a different opcode set, so it is treated structurally in §5, not as a fifth column here. [HIGH/OBSERVED]

TONGA (V1): engine-scoped TONGA_ISA_TPB_OPCODE (44 entries), byte = local \| (ENGINE<<5), in the separate TongaArchIsa package. Its compute-datapath bytes are inherited by v2 verbatim (§6.1); its control/SIM bytes are re-homed/dropped. See §5.

The presence matrix below is the union across v2..v5 = 172 distinct opcode names (including the 0xff sentinel). add/change/remove are marked inline where they occur at a generation boundary. [HIGH/OBSERVED]

Matrix dimensions: 172 opcode rows × 4 generation columns (V2/V3/V4-4+/V5). Per-column live totals (enum-body entries, sentinel included): V2 = 145, V3 = 150, V4 ≡ V4+ = 159, V5 = 165. [HIGH/OBSERVED]

CORRECTION (to DX-GEN-04 §2.1 / §5.3). DX-GEN-04 states JPEG_DECODE (0x81, +v3) "reuses a byte freed by a sunda removal" and calls 0x81 "the only reused freed byte." That is wrong. In the SUNDA opcode enum the sequence steps directly from DROPOUT = 0x7f (sunda common.h:215) to TRANSPOSE_BATCH_NORM_STATS2 = 0x82 (:216) — 0x80 and 0x81 are both holes, never occupied by any SUNDA opcode (and both stay holes in all four gens). So CAYMAN's JPEG_DECODE = 0x81 (cayman common.h:219) fills a pre-existing free byte, not a freed-and-reused one. The corrected invariant is stronger: v2→v3 performs zero byte-reuse — all 12 CAYMAN additions claim genuinely-free bytes, and all 7 retired SUNDA bytes are permanently abandoned. [HIGH/OBSERVED]

3. The per-step diff narrative

The lineage net-growth is 145 −(7) +(12) → 150 +(9) → 159 +(6) → 165. Exactly one step has removals (the v2→v3 re-baseline); v3→v4 and v4→v5 are strictly additive. [HIGH/OBSERVED]

3.1 SUNDA (v2) → CAYMAN (v3): 145 → 150 (+12 / −7)

The only step in the lineage with removals. It is a re-baselining: seven SUNDA-era ops are retired and twelve image/sparse/collective/codec ops land.

ADDED in CAYMAN (12) — all claim fresh (previously-unoccupied) bytes: LDTAGS 0x06, MATMUL_SPARSE 0x07 (PE low range); JPEG_DECODE 0x81, PSEUDO_JPEG_DECODE 0xd7 (codec); RANGE_SELECT 0xbc, DMA_TRANSPOSE 0xbd, GET_SEQUENCE_BOUNDS 0xbe, SB2SB_COLLECTIVE 0xbf (DMA/collective cluster); CONV_LUT_LOAD 0xe4, DVE_READ_INDICES 0xe9, DMA_GATHER_TRANSPOSE 0xf1, NONZERO_WITH_COUNT 0xf2 (high cluster). [HIGH/OBSERVED]

REMOVED at CAYMAN (7) — the only deletions in v2..v5, all abandoned (never reused v3..v5):

The freed band is {0x87,0x88,0x8a,0x8b,0x8c,0x8d,0x8f}; the sandwiched 0x8e BATCH_NORM_PARAM_LOAD2 survives all five gens (it is not in the retired band). The dual-pointer multi ops were already deprecated in SUNDA (// n); the BF16-fused arithmetic ops were live in v2 (// Y) and superseded at the v3 re-baseline by the generic TensorTensor/TensorScalar/TensorReduce arith ops carrying a dtype field. Anchors: sunda common.h:221–228. [HIGH/OBSERVED]

GOTCHA — forensic value of the abandoned band. Because the seven freed bytes are never reused in v3/v4/v5, a stray 0x87/0x88/0x8a/0x8b/0x8c/0x8d/0x8f in a v3+ instruction stream is a decode error or a v2 artifact, never a live v3+ instruction. Likewise 0x80/0x81 (with 0x81 taken by JPEG_DECODE only from v3) are SUNDA holes. A decoder validating against the wrong generation will mis-flag these. [HIGH/OBSERVED]

Second-source cross-check. The per-gen instruction_mapping.json struct2opcode (the wire-encodable subset) goes sunda 89 → cayman 99 (+10): of the 12 added, 10 carry a new encodable struct; the two pseudo-form additions (PSEUDO_JPEG_DECODE, plus the pseudo accounting) sit in the pseudo band. [HIGH/OBSERVED on the 89→99 totals; MED/INFERRED on the exact per-op struct accounting]

3.2 CAYMAN (v3) → MARIANA (v4): 150 → 159 (+9 / −0)

Pure-additive. MARIANA introduces the microscaling (MX) PE path, the sparsity-compress path, a second-gen RNG, the PE seed manager, ACTIVATE2, and an event-semaphore test op:

PE_MANAGE_SEED 0x08, LDWEIGHTS_MX 0x09, MATMUL_MX 0x0a (PE family, next free low bytes after SUNDA's 0x01–0x07); ACTIVATE2 0x25 (after the 0x21–0x24 ACT cluster); TEST_EVENT_SEM 0xb4 (between POLL_SEM 0xb3 and BRANCH_PREFETCH_HINT 0xb5); SPARSITY_COMPRESS 0xe0, SPARSITY_COMPRESS_TAG 0xe1, RAND2 0xe2, QUANTIZE_MX 0xe3 (the MX/sparse cluster, just below CONV_LUT_LOAD 0xe4). None of the seven v2→v3-freed bytes are touched. struct2opcode: cayman 99 → mariana 108 (+9) — the cleanest cross-source confirmation in the lineage (all nine are encodable, none pseudo). [HIGH/OBSERVED]

NOTE — QUANTIZE_MX 0xe3 binds DVE, not POOL. QUANTIZE_MX is the MX-quantize handler that runs on the DVE engine; it is absent from the POOL kernel-info tables (the named POOL QuantizeMx handler count drops 60→59 when it is moved out). The POOL MX surface is TENSOR_DEQUANTIZE 0x7b, not 0xe3. See the Cross-Gen Kernel-Info Matrix for the per-engine KIT binding. [HIGH/OBSERVED]

3.3 MARIANA (v4) → MAVERICK (v5): 159 → 165 (+6 / −0)

Pure-additive. MAVERICK adds the compact-control instruction, the DMA immediate/2D-copy fast forms, a second copy-memcpy, a multipass activate, and the INT_WIDE tensor ops:

ACTIVATE_MULTIPASS 0x26 (after MARIANA's ACTIVATE2 0x25); COMPACT_CONTROL_INST 0xb6, DMA_MEMCPY2 0xb9, DMA_IMMEDIATE 0xba (control/DMA cluster, around the existing 0xb5/0xb8/0xbb); TENSOR_TENSOR_INT_WIDE 0xf3, TENSOR_SCALAR_INT_WIDE 0xf4 (after MARIANA-era 0xf1/0xf2). struct2opcode: mariana 108 → maverick 114 (+6), all six encodable. Anchors: maverick common.h:173, 266, 268, 269, 320, 321. [HIGH/OBSERVED — header values OBSERVED; v5 interior semantics INFERRED]

GOTCHA — slot-policy is topical, not sequential. New ops do not simply take the next free byte globally; they take the next free byte in a topical cluster (PE family low, ACT 0x2x, control/DMA 0xb*, high cluster 0xe*/0xf*). This is why the additions are interleaved with invariant rows, and why the abandoned 0x8x band is never backfilled even though plenty of high-numbered bytes were free. The byte's high nibble is a soft engine-affinity hint, not a hard field (in v2+; in V1 it was a hard 3-bit field — §5.2). [MED/INFERRED]

4. The value-stability proof

Every opcode present in two generations carries the identical hex value in both. The pairwise value-mismatch join is empty for sunda↔cayman, cayman↔mariana, mariana↔maverick, and the full sunda↔maverick span (138 common opcodes, zero mismatches). The opcode byte is a write-once global namespace: new generations only append, removals only abandon, nothing is ever renumbered. [HIGH/OBSERVED]

A representative cross-range sample, value-identical across all of V2/V3/V4/V5 (verified directly from the four enum bodies):

The strong form is the sunda↔maverick (full-span) check: every one of the ~138 opcodes that survive from v2 to v5 keeps its v2 byte unchanged through three generation steps. This is the gen-invariance evidence that lets a reimplementer hard-code one opcode→handler table and parameterize only the presence set per generation (see The Gen-Invariance Thesis). And the invariant's first link is V1→V2 — the entire V1 compute-datapath byte namespace is inherited verbatim by SUNDA (§6.1). [HIGH/OBSERVED]

The union across v2..v5 is 172 distinct opcode names; the live set is 165 (the MAVERICK superset). 172 − 7 (the SUNDA-only retired ops) = 165. Every live op is present in MAVERICK, which is a strict superset of {CAYMAN ∪ MARIANA}. On the opcode axis the chain is SUNDA′ ⊊ CAYMAN ⊊ {MARIANA ≡ MARIANA_PLUS} ⊊ MAVERICK (where SUNDA′ = SUNDA minus the 7 retired). [HIGH/OBSERVED]

NOTE — opcode axis ≠ register/CSR axis. On the register/CSR axis the inclusion is SUNDA ⊂ {CAYMAN ≡ MARIANA ≡ MARIANA_PLUS} ⊂ MAVERICK (CAYMAN and MARIANA share a CSR schema). On the opcode axis CAYMAN ⊊ MARIANA strictly (MARIANA adds 9 opcodes). The ≡ between MARIANA and MARIANA_PLUS holds on both axes; the ≡ between CAYMAN and MARIANA holds only on the CSR axis. The two are not contradictory — they describe different axes. [HIGH/OBSERVED]

5. TONGA (V1): the pre-unified outlier

TONGA is NeuronCore-v1 — the first-generation Trn1/Inf1-era Tensor-Processing Block (TPB). It is outside the GPSIMD family: it has no coretype, no arch_id, no NCFW firmware image, no EXTISA blob, and zero runtime strings; it lives only in legacy ISA/register headers, and the SPIS Product-ID register places it one generation before SUNDA (Tonga 0x01 < Sunda 0x02 < Cayman 0x03). See The Codename → Generation Map for its identity placement. This section recovers what V1's ISA actually is. [HIGH/OBSERVED]

CORRECTION (to DX-GEN-04 §1.2 / §7). DX-GEN-04 reports TONGA as "register-block RTL only, no opcode enum," and marks the matrix TONGA column n/a. That reading is correct only for the RTL register tree arch-headers/tonga/ (cmake TongaArchHeaders, 215 files). It is incomplete: there is a second TONGA tree — arch-isa/aws_tonga_isa_* (cmake TongaArchIsa, copyright 2018) — which does ship a real typedef enum TONGA_ISA_TPB_OPCODE (44 entries) plus the 64-byte instruction-struct family and a separate SP sequencer ISA. V1 has a structured opcode/struct ISA; it simply uses an engine-scoped encoding rather than the unified flat namespace. The n/a is right for the flat axis but V1's bytes are mappable onto it (§6.1). [HIGH/OBSERVED]

5.1 The V1 NeuronCore geometry

TONGA_ISA_TPB_ARCH_CONSTS (arch-isa/tpb/aws_tonga_isa_tpb_common.h:13–37) — the classical fixed-function TPB: [HIGH/OBSERVED]

This is a PE-array + ACT + POOL over SB/PSUM. There is no 8-core Q7-POOL DSP cluster, no IRAM/DRAM Q7 apertures, no Vision-Q7 microarch — those are the v2+ GPSIMD additions (§7). [HIGH/OBSERVED]

5.2 The opcode construction rule: engine ‖ local

V1's opcode namespace is engine-scoped: the engine field is encoded in the opcode byte. From aws_tonga_isa_tpb_common.h:40–48, 65–67:

typedef enum TONGA_ISA_TPB_ENGINE {            // :40
    TONGA_ISA_TPB_ENGINE_PE      = 0x0,
    TONGA_ISA_TPB_ENGINE_ACT     = 0x1,
    TONGA_ISA_TPB_ENGINE_POOL    = 0x2,
    TONGA_ISA_TPB_ENGINE_ALL     = 0x3,
    TONGA_ISA_TPB_ENGINE_RT      = 0x6,        // instructions patched by the runtime
    TONGA_ISA_TPB_ENGINE_SIM     = 0x7,        // simulator-only, not in hardware
    TONGA_ISA_TPB_ENGINE_INVALID = 0xFF
} TONGA_ISA_TPB_ENGINE;

typedef enum TONGA_ISA_TPB_OPCODE {            // :65
    TONGA_ISA_TPB_OPCODE_LDWEIGHTS = 0x01 | (TONGA_ISA_TPB_ENGINE_PE   << 5),  // :67 ⇒ 0x01
    TONGA_ISA_TPB_OPCODE_POOL      = 0x05 | (TONGA_ISA_TPB_ENGINE_POOL << 5),  // :87 ⇒ 0x45
    TONGA_ISA_TPB_OPCODE_NOP       = 0x08 | (TONGA_ISA_TPB_ENGINE_ALL  << 5),  // :106 ⇒ 0x68
    /* ... */
} TONGA_ISA_PACKED TONGA_ISA_TPB_OPCODE;

The 1-byte opcode is [ engine:3 (bits 7..5) ‖ local_op:5 (bits 4..0) ]. DECODE: engine = opcode >> 5, local = opcode & 0x1f. The byte carries structure. The unified v2+ NEURON_ISA_TPB_OPCODE inherits the byte values but drops the field semantics — the engine becomes an out-of-band dispatch property, not a bitfield, and the byte is frozen as a flat write-once key. This is the genesis of the "stable global opcode byte" that §4 proves invariant v2..v5. [HIGH/OBSERVED]

A further structured field survives inside the encoding: within the POOL engine the 0x10 local bit is the arith/bitvec discriminator. E.g. TENSOR_TENSOR_ARITH_OP is local 0x01 → byte 0x41, while TENSOR_TENSOR_BITVEC_OP is local 0x11 → byte 0x51 (aws_tonga_isa_tpb_common.h:75–76). The unified enum preserves both bytes (0x41, 0x51) but no longer documents the bit as a field. [HIGH/OBSERVED]

5.3 The 44-entry V1 roster (byte = local ‖ engine<<5)

(RT ops are runtime-patched placeholders, "never executed by H/W"; SIM ops are simulator-only.) [HIGH/OBSERVED]

5.4 The 64-byte TONGA_ISA_TPB instruction-struct family

V1's instruction format is a 64-byte fixed record (the v2+ model is a flat operand-struct keyed off the opaque opcode; V1's is engine-scoped and self-validating per instruction). The umbrella aws_tonga_isa_tpb.h pulls in 39 struct headers: 25 compute/ALL (always), 8 RT-pseudo, 5 SIM-only (under #ifdef SIM_ONLY). Every compute struct STATIC_ASSERTs to INST_NBYTES = 64, and each ships a tonga_isa_tpb_<op>_check_validity() — this is a wire-format-exact host encoder/validator library, not an internal model. [HIGH/OBSERVED]

The shared 8-byte prefix (every TPB instruction):

TONGA_ISA_TPB_INST_HEADER (4B, common.h:214–226):
    +0 opcode(1)  +1 inst_word_len(1)  +2 debug_cmd(1)  +3 debug_hint(1)
TONGA_ISA_TPB_INST_EVENTS (4B, common.h:229–235):
    +0 wait_event_mode(1)  +1 wait_event_idx(1)  +2 set_event_mode(1)  +3 set_event_idx(1)

The embedded per-instruction wait-on-event / set-on-done synchronization is part of every compute instruction — V1's dataflow sync model. The memory-access patterns MEM_ACCESS_1D/2D/3D/4D are 8/12/16/20 bytes ({ tpb_addr start_addr; int16 step_elem[N]; uint16 num_elem[N]; }, i.e. 4 + 4N; common.h:246–269). [HIGH/OBSERVED]

Exemplar — TONGA_ISA_TPB_COPY_INST (64 B, opcode 0x46, POOL engine) (aws_tonga_isa_tpb_copy.h:14–49):

Header doc: "Copy … done through the pooling engine."

Exemplar — TONGA_ISA_TPB_MATMUL_INST (64 B, opcode 0x02, PE engine) (aws_tonga_isa_tpb_matmul.h:96–197): inst_header(4) + inst_events(4) + in_dtype(1) + perf_opt(1: PE_PERF_OPT_MODE) + psum_accumulate_flags(1) + ifmap_replication_{resolution,shift_amnt,num_rows}(3) + quant_offset union(2) + src/dst MEM_ACCESS_3D(16+16) + num_active_rows(1) + num_active_cols(1) + reserved[14] == 64. Math (doc :10): Psum_out = (Xq−Xqz)·(Wq−Wqz) + Psum_in over the 128×64 array — native asymmetric UINT8/UINT16 quantization with zero-points; UINT8 DOUBLE_ROW/COLUMN/PIXEL 2×/cycle perf modes. Note MatMul uses MEM_ACCESS_3D (16 B) while COPY/POOL use MEM_ACCESS_4D (20 B). [HIGH/OBSERVED]

5.5 The SP sequencer ISA (separate, 8-byte, opcode:7 + phase:1)

The Sequencer (SP) is a separate instruction set with its own encoding (8-byte instructions: SP_INST_NBYTES = 8, _NWORDS = 2). The instruction header is a bitfield { opcode:7; p:1 } (aws_tonga_isa_sp_common.h:45–48), masked OPCODE_MASK = 0x7F / P_MASK = 0x80 (aws_tonga_isa_sp_instruction_fields.h). The 16 SP opcodes (aws_tonga_isa_sp_instruction_opcodes.h:9–24): [HIGH/OBSERVED]

The SP is V1's control-flow / event / counter / branch engine — the on-chip program sequencer that drives the TPB engines, with full 64-bit wire-field diagrams in sp_isa.txt. [HIGH/OBSERVED]

5.6 The V1 dtype family

TONGA_ISA_TPB_DTYPE (1-byte packed, common.h:51–62) is the distinct name family (TONGA_ISA_TPB_DTYPE_* vs the unified NEURON_ISA_TPB_DTYPE_*) that marks V1 as the older arch — but the byte-width-encoded values carry forward into the v2+ dtype family: INVALID 0x0, UINT8 0x3, UINT16 0x5, BFLOAT16 0x6, FP16 0x7, INT32 0x8, FP32 0xA, INT64 0xC. [HIGH/OBSERVED]

6. The V1→V2 transition

Computed by joining V1's TONGA_ISA_TPB_OPCODE bytes (engine<<5 | local) against the unified SUNDA NEURON_ISA_TPB_OPCODE. Of V1's 44 entries, 30 share a name with SUNDA; of those, 26 are byte-identical and 4 are re-homed; 14 are V1-only. [HIGH/OBSERVED]

6.1 Carried forward byte-identical — the compute datapath

The entire V1 compute-datapath byte namespace is inherited verbatim by V2: [HIGH/OBSERVED]

• PE: LDWEIGHTS 0x01, MATMUL 0x02.
• ACT: ACTIVATE 0x21, ACTIVATE_QUANTIZE 0x22.
• POOL (22 ops): the whole 0x41–0x5e block — TENSOR_TENSOR/REDUCE/SCALAR/ SCALAR_PTR/CUMULATIVE (arith 0x41–0x4e + bitvec 0x51–0x5e), POOL 0x45, COPY 0x46, CAST 0x47, RECIPROCAL 0x48, MEMSET 0x49, REG_LOAD/STORE/SHUFFLE 0x4a–0x4c, RNG 0x4d.
• PSEUDO: PSEUDO_SEMAPHORE_SET 0xc5, PSEUDO_LOAD_ACT_FUNC_SET 0xc6.
• SENTINEL: INVALID 0xff.

This is the birth of the write-once opcode byte that §4 proves invariant v2..v5 — the invariant's first link is established at V1→V2. [HIGH/OBSERVED]

6.2 Re-homed — the control/sync block moves out of ENGINE_ALL

V1 put cross-engine control ops in the ENGINE_ALL (3<<5 = 0x60) band; V2 rebuilt control/sync into the 0xa0–0xab cluster and replaced the engine-scoped event ops with the unified EVENT_SEMAPHORE family: [HIGH/OBSERVED]

The 4 re-homed shared ops are {NOP, WRITE, NOTIFY, PSEUDO_READ_VAR_ADDR}; the engine-scoped EVENT/SEM ops are replaced (not moved) by the unified EVENT_SEMAPHORE model. The compute surface V2 kept; the control surface it redesigned. [HIGH/OBSERVED]

6.3 Dropped at V2 — V1-specific bands

The engine-scoped EVENT_{WAIT,SET,CLEAR} / SEMAPHORE_TEST_AND_SET (replaced); the RT-pseudo PSEUDO_DMA_{TRIGGER 0xc1, REARM 0xc2, BARRIER 0xc3, MEMCPY 0xc4} + PSEUDO_SET_TRANSFER_ADJUST 0xc7 (V2 keeps the byte band 0xc1–0xc4 for PSEUDO_DMATRIGGER/DMAREARM/DMABARRIER/DMAMEMCPY_FULL_IND — same bytes, renamed and reorganized into the unified 0xc1–0xdf PSEUDO band); and the entire SIM_* band (SIM_DMA_COPY 0xeb … SIM_MEMCPY 0xef, ENGINE_SIM 7<<5) — simulator-only ops that V2 does not carry into the runtime enum. [HIGH/OBSERVED]

6.4 What the transition means

V2 SUNDA is V1 TONGA with (a) the engine-field encoding flattened to an opaque byte, (b) the control/sync block redesigned around EVENT_SEMAPHORE, (c) the SIM band dropped, and (d) the POOL engine replaced by the programmable NX_POOL GPSIMD DSP (§7) — while keeping the producer-side opcode bytes. The compute-datapath bytes are the conserved core; everything else is re-homed. [HIGH/OBSERVED on (a)(b)(c); MED/INFERRED on (d)'s "producer-side bytes kept across a silicon swap" framing — the cleanest reading of why POOL 0x45 survives a complete engine replacement.]

7. The GPSIMD-genesis answer

Is there a Vision-Q7 GPSIMD POOL engine in V1? NO. GPSIMD — the programmable Cadence Tensilica Vision-Q7 NX_POOL cluster running the Vision-Q7 device ISA, with custom-op image download, DVE, and IRAM/DRAM apertures — is a SUNDA-onward (v2+) feature. V1 predates it. [HIGH/OBSERVED]

Three independent shipped-artifact lines of evidence:

1. V1's POOL is fixed-function, not programmable. aws_tonga_isa_tpb_pool.h defines POOL_INST (0x45) as a hardwired planar reduction: a 2-value pool_func enum (MAXPOOL = 0x01, AVGPOOL = 0x02; pool.h:54–58), a pool_dim (TENSOR_SUBDIM), and a pre-baked pool_scale float (= 1/(R·S) for averaging). The doc states: "applies a planar function … Supported pooling functions are MaxPool and AveragePool. All pooling calculations are performed in FP32" (pool.h:6–11). No downloadable kernel, no custom-op image, no instruction RAM, no Q7 core — a fixed pool_func selector, not a kernel. The 64-byte POOL_INST layout: inst_header(4)+inst_events(4)+pool_func(1)+in_dtype(1)+ out_dtype(1)+pool_dim(1)+pool_scale(4)+src/dst MEM_ACCESS_4D(20+20)+ num_active_channels(1)+reserved[7] == 64 (pool.h:61–137). [HIGH/OBSERVED]

2. Zero GPSIMD / Vision-Q7 / DVE / custom-op footprint in the V1 ISA. A case-insensitive scan of the entire TONGA arch-isa/ tree for gpsimd|dve|q7|ncore2gp|vision.?q7|custom.?op returns 0 matches across every term. The V1 ISA has no concept of a programmable DSP cluster, no DVE, no custom-op descriptor, no Q7 apertures. [HIGH/OBSERVED]

3. Every runtime core is an NX_POOL core — and there is no TONGA core. The shipped runtime container exposes exactly NRTUCODE_CORE_{SUNDA,CAYMAN,MARIANA,MARIANA_PLUS,MAVERICK}_NX_POOL and the per-gen image tables sunda_libs … maverick_libs. Every v2+ core carries the _NX_POOL suffix (NX = the Vision-Q7 GPSIMD POOL engine). There is no TONGA core, no tonga_libs, and zero "tonga" strings in the runtime lib. The TONGA ISA package is compile-time-only (a host-side encoder/validator + cmake export), with zero runtime identity. [HIGH/OBSERVED]

The genesis, stated. The POOL name and the producer-side opcode byte 0x45 are conserved across the V1→V2 boundary, but the thing behind them is swapped: V1 POOL = fixed-function MaxPool/AvgPool/vector engine; V2+ NX_POOL = the programmable Vision-Q7 GPSIMD cluster (the device DSP this entire book documents). The producer-side TPB opcode byte is kept so the host emitter and the NEFF wire format stay stable, but the consumer changed from fixed silicon to a downloadable-kernel DSP. GPSIMD is a v2+ feature; V1 is the pre-GPSIMD, fixed-function-POOL baseline. [HIGH/OBSERVED on the three evidence lines; MED/INFERRED on the silicon-swap-behind-conserved-opcode framing]

8. Is MARIANA_PLUS a distinct ISA? — No

MARIANA_PLUS (v4+, arch_id 28 / coretype 29) is NOT a distinct ISA. It is a recompile + flag-refresh of MARIANA (v4, arch_id 20 / coretype 21). It executes the MARIANA 159-opcode table verbatim — same key-set, same values, same KIT keys. [HIGH/OBSERVED]

The resolution rests on three independent shipped-artifact lines: [HIGH/OBSERVED]

1. No distinct ISA opcode table. There is no neuron_mariana_plus_arch_isa/ directory; the register tree arch-headers/mariana_plus/ contains no NEURON_ISA_TPB_OPCODE (it is CSR/register-block only). MARIANA_PLUS therefore runs the MARIANA 159-opcode enum, which is why the matrix V4/4+ column is a single shared column.
2. No version-axis point. No tree has a V4_PLUS / V4P enumerator in NEURON_ISA_TPB_NEURON_CORE_VERSION (§1.1). The runtime distinguishes v4 from v4+ via a string selector (NEURON_RT_DBG_V4_PLUS=0/1, formerly the compile flag NRTUCODE_MPLUS_ON_MARIANA), not an ISA version number.
3. It is a first-class firmware core with byte-stable ISA. MARIANA_PLUS is a peer NRTUCODE_CORE_MARIANA_PLUS_NX_POOL core with its own image set and register-map dir — but every getter name, dispatch handler, opcode value, PROF table, reset/boot geometry, and EXTISA Q7 kernel container is byte-stable or byte-identical v4 ↔ v4+. The only functional addition is one gen-wide firmware feature — the DGE reshape fast-path (a microcode/register behavior, not an opcode-table change, which is precisely why no new opcode appears).

On the generation axis: v2 SUNDA < v3 CAYMAN < v4 MARIANA ≤ v4+ MARIANA_PLUS < v5 MAVERICK, where ≤ means ISA-equal, firmware/register-revised. For the full byte-by-byte v4↔v4+ delta (the DGE fast-path, the arch_id/coretype algebra, the register-map refresh), see MARIANA_PLUS (v4+) Generation Delta. [HIGH/OBSERVED on the ISA-equal relation; the product-revision label is MED/INFERRED]

9. Corrections to the backing reports

Two in-place corrections were established from the bytes this pass. [HIGH/OBSERVED]

• 0x81 is a pre-existing hole, not a reused freed byte. DX-GEN-04 §2.1/§5.3 call JPEG_DECODE 0x81 "the only reused freed byte." The SUNDA opcode enum jumps DROPOUT 0x7f → TRANSPOSE_BATCH_NORM_STATS2 0x82 (sunda common.h:215–216); 0x80 and 0x81 are never occupied in any generation before CAYMAN fills 0x81 (and 0x80 stays a hole in all four). So v2→v3 performs zero byte-reuse. (Recorded in the matrix and §3.1.)
• TONGA V1 has a real opcode enum. DX-GEN-04 §1.2/§7 mark TONGA as "register-block only, no opcode enum." That holds for arch-headers/tonga/ (the RTL tree) but not for arch-isa/aws_tonga_isa_* (cmake TongaArchIsa), which ships a real engine-scoped TONGA_ISA_TPB_OPCODE (44 entries) + 64-byte struct family + SP ISA. (Recorded in the §5 CORRECTION callout.)

10. Confidence ledger

Related pages: Cross-Gen Kernel-Info Matrix · Arch-ISA Header Diff · Opcode Catalog Ledger · Gen-Invariance Thesis · Codename → Generation Map · MARIANA_PLUS (v4+) Delta · Confidence & Walls Model.