The 64-Byte Instruction Record Format

All addresses, offsets, 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, ELF64, not stripped, DWARF present, build-id 8bb57aba…). .text VMA equals file offset for the cited ranges. Struct offsets are read from DWARF (structures.json), enum values from DWARF (enums.json), and the encode/decode bodies from the disassembly. Other versions will differ.

Abstract

Every instruction the Neuron sequencer fetches — on every engine, for every NeuronCore generation V2/V3/V4 — is a fixed 64-byte record. TPB_INST_NBYTES is 0x40; the instruction stream is a flat array of these records, walked in +64 strides with no length prefix, no alignment padding, and no variable-width encoding. This is the runtime-side analog of a classic fixed-width VLIW/RISC instruction word: where a conventional ISA packs a 32- or 64-bit word, the TPB ISA packs a 64-byte word, because each record must carry a full tensor descriptor (3-D start/step/count), an inter-engine synchronization event tuple, immediates, and a multi-byte opcode payload in one atomic fetch. Contrast the libtpu BarnaCore 32-byte bundle, which is a true bit-packed VLIW filled field-by-field by a BitCopy(dst_bitoff, …) primitive: the TPB record is byte-granular and struct-shaped — a tagged union whose first four bytes are a NEURON_ISA_TPB_HEADER and whose remaining 60 bytes are one of 118 distinct C structs selected by the opcode byte. There is no bit-packer; fields are plain uint8/uint16/uint32 members at fixed byte offsets, written by the compiler and read back by memcpy-shaped accessors.

The record is a discriminated union. Byte [0] is the opcode; the remaining 63 bytes are reinterpreted as the matching NEURON_ISA_TPB_<OP>_STRUCT (e.g. CTRL_AL_STRUCT for an ALU op, DMA_DIRECT2D_STRUCT for a 2-D DMA, PSEUDO_FUNCTION_CALL_STRUCT for a call). All 118 arms share a 12-byte preamble — a 4-byte header (opcode, inst_word_len, debug_cmd, debug_hint) and an 8-byte EVENTS tuple (wait_mode/wait_idx/update_mode/update_idx/semaphore_value) — and then diverge into per-opcode payload from byte [12] onward. The opcode space splits cleanly: 0x01..0xFF real (hardware-executed) opcodes, and a 0xC1..0xDF pseudo-opcode band that the host-side lowering pass (Pseudo-Instruction Lowering) rewrites into real records before the record ever reaches the device. Two records never leave the host: PSEUDO_FUNCTION_BEGIN (0xD1) splits the stream into functions, and PSEUDO_INST (0xDF) with sub-opcode 1 is the EXIT_EXECUTION sentinel that terminates a program.

This page documents three artifacts a reimplementer must reproduce: (1) the 64-byte union layout — the shared 12-byte preamble plus the per-opcode payload, pinned to byte offsets from DWARF; (2) the opcode taxonomy — the 70 real + 50 pseudo opcode values and the field-group dimensions that vary across opcode classes, given as axes rather than a 118-row dump; and (3) the instruction-block I/O algorithms — the three ib_* workhorses (ib_read_write, ib_has_pseudo_exit, ib_final_inst_validity_check) that copy, scan, and validate a record stream as it is staged from host into device HBM. The per-opcode legality rules are owned by the ISA Validator Architecture; the per-arch deltas by the Per-Arch Validators page; the host-to-real lowering of the pseudo band by Pseudo-Instruction Lowering.

For reimplementation, the contract is:

• The fixed-width model: TPB_INST_NBYTES = 0x40; the stream is record[i] at byte i*64; every walker uses a +0x40 stride and byte[0] as the discriminant. There is no framing byte and no length field in the stream itself.
• The shared 12-byte preamble: NEURON_ISA_TPB_HEADER (4 B at +0) + NEURON_ISA_TPB_EVENTS (8 B at +4), identical across all 118 union arms.
• The opcode→struct dispatch: opcode (byte +0) selects one of 118 NEURON_ISA_TPB_<OP>_STRUCT reinterpretations of bytes +12..+63.
• The pseudo band: 0xC1..0xDF (plus 0xA9 COMPARE_BRANCH, 0xB5/0xDD branch-hint) is host-only; 0xD1 FUNCTION_BEGIN and 0xDF/sub-1 EXIT are stream sentinels.
• The ib_* I/O contract: a record stream is copied to/from device dmem segments by ib_read_write (chunked, segment-split via ib_translate_addr), scanned for the EXIT sentinel by ib_has_pseudo_exit, and final-validated by ib_final_inst_validity_check.

1. The Record Layout

Purpose

The 64-byte record is the single physical unit the entire model-load and execute path manipulates. The NEFF carries it, kelf2kbin emits it, the lowering pass rewrites it, the validator checks it, ib_read_write stages it into HBM, and the sequencer fetches it. Reading every downstream offset table requires first internalizing the shared shape: bytes [0..11] are opcode-independent; bytes [12..63] are opcode-dependent. A reimplementer who treats the record as opaque bytes will mis-place every field; one who internalizes the 12-byte preamble and the opcode→struct table can decode any record.

Encoding

The record is a C union whose every arm begins with the same 12-byte preamble. The preamble is two structs:

NEURON_ISA_TPB record  (64 bytes; TPB_INST_NBYTES = 0x40)

byte: 0      1            2          3        4 ........... 11   12 ......................... 63
     +------+------------+----------+--------+------------------+------------------------------+
     |opcode|inst_word_len|debug_cmd|dbg_hint|   EVENTS (8 B)   |   per-opcode payload (52 B)  |
     | [0]  |    [1]      |   [2]   |  [3]   |   wait/update     |   reinterpreted by opcode    |
     +------+------------+----------+--------+------------------+------------------------------+
       \________ HEADER (4 B @+0) ________/  \__ EVENTS @+4 __/  \____ union arm @+12 _________/

NEURON_ISA_TPB_HEADER (4 bytes @+0)

NEURON_ISA_TPB_EVENTS (8 bytes @+4)

The synchronization tuple every instruction carries. It is how the TPB ISA expresses inter-engine ordering inline in each record rather than via separate barrier instructions — a wait predicate and an update action both ride in the 8 bytes.

NOTE — EVENTS has an extended 12-byte sibling NEURON_ISA_TPB_EVENTS_EXTENDED that splits semaphore_value into separate sem_wait_value (+4) and sem_update_value (+8). The 8-byte form is the one embedded at +4 in the standard 64-byte record; the extended form appears in records that need independent wait/update immediates. A reimplementer decoding a record reads the 8-byte form unless the opcode struct declares the extended one.

The per-opcode payload (+12..+63)

Bytes +12 onward are reinterpreted by the opcode. There are 118 distinct 64-byte arm structs in DWARF; their first 12 bytes are always HEADER+EVENTS, so they differ only in the 52-byte payload. The payload regions recur across families:

Function Map

GOTCHA — inst_word_len (byte +1) is not the stream stride. Every record occupies exactly 64 bytes in the stream and every walker (ib_has_pseudo_exit, the lowering splitter, the validator) advances by 0x40. inst_word_len is a sequencer-fetch hint, not a record size; a reimplementer that strides by inst_word_len will desynchronize the stream immediately.

2. The Opcode Taxonomy

Purpose

The opcode byte at [0] is the discriminant for the entire union. There are 120 defined opcode values across two bands: a real band the hardware executes, and a pseudo band the host lowers away. Rather than dump all 118 arm structs, this section gives the opcode value space and the field-group dimensions that vary across opcode classes — the axes a reimplementer reconstructs the per-opcode layout from.

Encoding

The real opcodes cluster by engine/datapath in value ranges; the pseudo opcodes occupy a contiguous 0xC1..0xDF band plus three control-flow stragglers.

Real opcode value clusters (70 values, abbreviated by range)

CORRECTION (ISA-IR-2) — an earlier revision of this cluster table placed JPEG_DECODE=0xFF in the 0xE0..0xFF row alongside INVALID=0xFF. That value was wrong. The authoritative DWARF NEURON_ISA_TPB_OPCODE enum (1-byte underlying type) gives JPEG_DECODE=0x81, its host-lowered CC-top form PSEUDO_JPEG_DECODE=0xD7, and 0xFF=INVALID (the terminal sentinel only). 0xFF is not JPEG decode; a reimplementer keying the JPEG-decode path off 0xFF would never match the real opcode and would collide with the INVALID sentinel. This matches the ISA overview sibling page (PSEUDO_JPEG_DECODE 0xD7).

Pseudo opcode band (host-only; lowered before device fetch)

PSEUDO_INST (0xDF) is an extension prefix: byte +12 is a NEURON_ISA_TPB_PSEUDO_OPCODE (INVALID=0, PSEUDO_EXIT_EXECUTION=1, PSEUDO_LIBRARY_RELOAD_INDEX=2) that sub-dispatches the record. This two-level encoding is why ib_has_pseudo_exit checks both byte[0]==0xDF and byte[12]==1.

CORRECTION (TDRV-CORE-15) — an earlier note labeled opcode 0xDF simply "exit". It is the PSEUDO_INST extension prefix; the EXIT sentinel is specifically PSEUDO_INST with pseudo_opcode==1 (PSEUDO_EXIT_EXECUTION, the NEURON_ISA_TPB_PSEUDO_EXIT_EXECUTION_STRUCT whose pseudo_opcode field sits at +12). A scanner that keys on byte[0]==0xDF alone will false-positive on the LIBRARY_RELOAD_INDEX sub-form (pseudo_opcode==2).

The field-group dimension table

The 118 arm structs are not 118 independent layouts — they are combinations of a small set of payload field-groups. A reimplementer reconstructs any arm from these axes plus the shared preamble:

Worked arm layouts

Four representative arms, pinned to DWARF offsets, show the dimension combinations:

CTRL_AL (ALU op; opcode in CTRL band)               DMA_DIRECT2D (opcode 0xB.. 2-D DMA)
 +0  HEADER (4)                                       +0  HEADER (4)
 +4  EVENTS (8)                                       +4  EVENTS (8)
 +12 op / dtype / src_datasrc / dst_dtype (4×1)       +12 dma_configs / semaphore / sem_increment / compute_op
 +16 src0_lo/hi src1_lo/hi dst_lo/hi (6× REG_NUM)     +16 src_start_addr (ADDR8, 8)
 +22 reserved1 (2)                                    +24 src_step_elem (int32[2]) / +32 src_num_elem (u16[2])
 +24 immediate (ALU_IMMEDIATE, 8)                     +36 src_elem_size / +38 src_bound_reg / +39 dst_bound_reg
 +32 reserved2 (32)                                   +40 dst_start_addr (ADDR8, 8) / +48 dst_step_elem ...
                                                      +62 in_dtype / +63 out_dtype

S3D3_TS (tensor-scalar; two TENSOR3D + 2 imms)        PSEUDO_FUNCTION_CALL (opcode 0xD3, host-only)
 +0  HEADER (4) / +4 EVENTS (8)                        +0  HEADER (4) / +4 EVENTS (8)
 +12 accumulator_cmd / reserved0 (3)                   +12 function_name (int8[36])
 +16 src_mem_pattern (TENSOR3D, 16)                    +48 args_table_var_id (uint32)
 +32 in_dtype / out_dtype / num_active_channels        +52 reserved0 (12)
 +35 imm0_src / op0 / op1 / reverse_operands / imm1_src
 +40 imm0 (4) / +44 imm1 (4)                          PSEUDO_FUNCTION_BEGIN (opcode 0xD1, host-only)
 +48 dst_mem_pattern (TENSOR3D, 16)                    +12 function_name (int8[36])
                                                       +48 return_reset_semaphores / +49 return_addr_reg_lo
                                                       +50 return_addr_reg_hi / +51 reserved0 (13)

QUIRK — the PSEUDO_FUNCTION_BEGIN record stores return_addr_reg_lo at +49 and return_addr_reg_hi at +50, and return_reset_semaphores at +48. The lowering pass reads exactly these three bytes to build the function's call/return frame (itf_identify_functions); a reimplementer of the lowering pass must read them at these offsets, not from a separate metadata table. The 36-byte function_name at +12 is the hash-table key.

3. The Operand Sub-Structures

Purpose

The payload field-groups are themselves small structs reused across arms. Pinning their byte layout lets a reimplementer decode any tensor descriptor, address operand, or register field without re-deriving it per opcode.

Encoding

The ADDR4/ADDR8 marker discipline

An address operand is either an immediate (a 29-bit ISA byte address in the low bits of the uint32/uint64) or register-sourced, distinguished by the marker byte at the top of the operand:

• ADDR4.marker & 0xA0 == 0x80 → register addressing (the regnum field is a live register selector).
• ADDR4.marker & 0x60 == 0x20 → indirect addressing (a data_addr field participates).
• Immediate form: addr_immediate & 0x1FFFFFFF is the 29-bit address; quadrant bits & 0x1E000000 select the SBUF/PSUM aperture.

NOTE — the 29-bit ISA address space is encoding address, not physical SRAM address. Its quadrant apertures (q0 0..0x37FFF, q1 base 0x800000, q2 0x1000000, q3 0x1800000, each 0x38000 wide; PSUM 0x2000000..0x2203FFF) are decoded by the validator's addresses_in_same_sbuf_quadrant / tpb_addr_active_channels predicates. The ISA-address-to-physical mapping is done elsewhere (the per-arch device geometry); these apertures are partition-relative and arch-independent at the record level. See Per-Arch Validators.

4. Instruction-Block I/O: ib_read_write

Purpose

Once the host has built a record stream for an engine, the stream must be copied into device-resident dmem (HBM/TDRAM) so the sequencer can fetch it — and copied back out for inspection. ib_read_write is the bidirectional workhorse: one function handles both directions, picking the dmem copy primitive by a read/write flag and splitting each transfer at dmem-segment boundaries via ib_translate_addr.

Entry Point

kbl_instruction_stream_read / kbl_instruction_stream_write
  └─ ib_read_write (0x323cc0, 284 B)
       ├─ ib_translate_addr   (0x323ca0, neighbour cell)  ── map linear offset → (dmem, off)
       ├─ dmem_buf_copyout    (read direction)            ── device → host
       └─ dmem_buf_copyin     (write direction)           ── host → device

Algorithm

// Models ib_read_write @0x323cc0 (284 B). r8b = is_read flag.
// eng_ib_addrs is an ib_addrs_one_eng_t (56 B): base dmem + 3 ib_code_range segments.
function ib_read_write(eng_ib_addrs, offset, buffer, size, is_read):
    copy_func = is_read ? dmem_buf_copyout      // 0x323ce4: read = device→host
                        : dmem_buf_copyin        // 0x323db8: write = host→device
    if size == 0: return NRT_SUCCESS             // 0x323cee
    cursor = 0
    while cursor < size:                         // 0x323d50 loop head
        found_dmem0 = found_dmem1 = 0
        dmem_offset = 0
        // map the current linear stream offset to a (dmem segment, in-segment offset)
        if ib_translate_addr(eng_ib_addrs, offset,           // 0x323d81
                             &found_dmem0, &dmem_offset, &found_dmem1) != 0:
            return NRT_FAILURE                   // 0x323dc8 → return 1
        dmem = found_dmem0 ? found_dmem0 : found_dmem1        // 0x323d8a/0x323d94
        seg_end  = found_dmem0 ? dmem_offset : 0
        // clamp this chunk to the segment boundary (cmovnb at 0x323d3b)
        seg_size = dmem.size - dmem_offset                   // 0x323d2a/0x323d31
        chunk    = min(seg_size, size - cursor)              // cmovnb: don't overrun size
        copy_func(dmem, buffer + cursor, dmem_offset, chunk) // 0x323d42
        cursor += chunk                          // 0x323d45
        offset += chunk                          // 0x323d48
    return NRT_SUCCESS                            // 0x323da4 → return 0

Encoding

ib_read_write operates on whole bytes, not records — but because the stream is a flat 64-byte array, any byte range it copies is a clean run of records when offset and size are multiples of 0x40 (which the callers ensure). The function's only structural knowledge of the record is indirect: ib_translate_addr maps a linear offset into one of three dmem segments (one_offset, model_switch_offset, core_offset in ib_addrs_one_eng_t @+0x08/+0x18/+0x28), so a record that straddles a segment boundary is copied in two chunks. The cmovnb at 0x323d3b is the boundary clamp.

Function Map

Considerations

The direction flag (r8b) selecting between two lea'd function pointers (0x323ce4 read, 0x323db8 write) is the entire read/write fork — there is no other behavioral difference between the two directions. A reimplementation can collapse to one loop with a function-pointer parameter, exactly as the binary does.

5. Instruction-Block I/O: ib_has_pseudo_exit

Purpose

A record stream must be checked for a program-terminating EXIT sentinel before it is staged (a stream missing its EXIT would run off the end). ib_has_pseudo_exit is the scanner: it walks the stream in 64-byte strides looking for the PSEUDO_INST/EXIT two-byte signature.

Entry Point

kbl_model_add
  └─ ib_has_pseudo_exit (0x323df0, 51 B)   ── scan instr_buf for the EXIT sentinel

Algorithm

// Models ib_has_pseudo_exit @0x323df0 (51 B). Verbatim 64-byte stride scan.
// instr_buf = base pointer (rdi); size = byte length (rsi/edx).
function ib_has_pseudo_exit(instr_buf, size):
    if size == 0: return false                   // 0x323df2: empty stream
    offset = 0
    while offset < size:                         // 0x323e04/0x323e07 (jnb)
        // record discriminant: opcode byte at the record base
        if instr_buf[offset + 0x00] == 0xDF      // 0x323e09: PSEUDO_INST prefix
           && instr_buf[offset + 0x0C] == 1:     // 0x323e0f: pseudo_opcode == PSEUDO_EXIT_EXECUTION
            return true                          // 0x323e16
        offset += 0x40                           // 0x323e00: TPB_INST_NBYTES stride
    return false                                 // 0x323e20

Encoding

This 51-byte function is the most compact independent proof of three record facts at once: (1) the stride is exactly 0x40 (add offset, 40h at 0x323e00); (2) the opcode is byte [0] (cmp byte ptr [instr_buf+offset], 0DFh at 0x323e09); and (3) the PSEUDO_INST sub-opcode is byte [+0x0C], i.e. +12 — the first payload byte past the 12-byte preamble (cmp byte ptr [instr_buf+offset+0Ch], 1 at 0x323e0f). The constant 1 is NEURON_ISA_TPB_PSEUDO_OPCODE_PSEUDO_EXIT_EXECUTION.

Function Map

GOTCHA — the sentinel check is a logical AND of two bytes 12 apart. Checking only byte[0]==0xDF matches every PSEUDO_INST record, including the LIBRARY_RELOAD_INDEX sub-form (pseudo_opcode==2). A reimplementer must check byte[12]==1 too, or a library-reload record will be mistaken for the program terminator.

6. Instruction-Block I/O: ib_final_inst_validity_check

Purpose

Before a staged record stream is committed, its final record is validated against per-arch rules (the last record must be a legal terminator for that engine on that NeuronCore generation). ib_final_inst_validity_check is the dispatch shim into the arch-specific validator.

Entry Point

sequencer_setup_instr
  └─ ib_final_inst_validity_check (0x323cb0, 5 B jmp thunk)
       └─ tdrv_arch_instr_block_final_inst_validity_check (arch vtable)
            └─ is_valid_<terminator> → is_valid_enum / type_size_check / ...

Algorithm

// Models ib_final_inst_validity_check @0x323cb0 — a 5-byte tail-call thunk.
function ib_final_inst_validity_check(kbin, eng_type, final_inst):
    // 0x323cb0: e9 9b 70 fe ff  → jmp tdrv_arch_instr_block_final_inst_validity_check
    return tdrv_arch_instr_block_final_inst_validity_check(kbin, eng_type, final_inst);
    // the arch vtable picks the SUNDA/CAYMAN/MARIANA validator and runs the leaf predicates

Encoding

The thunk is literally a single jmp (0x323cb0: e9 9b 70 fe ff). It exists so the call site (sequencer_setup_instr) has a stable, arch-neutral entry symbol while the actual validation routes through tdrv_arch_instr_block_final_inst_validity_check and its per-arch vtable. The leaf predicates it eventually reaches — is_valid_enum (the opcode/operand legality oracle, with the four OPCODE bitmasks 0x400082F07FBFFFFF / 0x024BC3FFFFDF3F7F / 0x7FEFFFFFE08F7FFF / 0x1001E0000003E), is_valid_dtype, type_size_check, and the address-quadrant checks — are documented on the ISA Validator Architecture page.

Function Map

NOTE — the validator primitives appear at multiple addresses (0x27b650 and 0x324050 for is_valid_enum; 0x27b9d0 and 0x324460 for type_size_check) — these are arch-specialized copies (SUNDA/CAYMAN/MARIANA), all decoding the same record format. The record layout on this page is arch-neutral; only the legality of specific field values varies per arch. See Per-Arch Validators.

Related Components

Cross-References

• Overview: the TPB Engine Instruction Model — where the 64-byte record sits in the engine pipeline
• Pseudo-Instruction Lowering (SP / TopSP Builders) — how the 0xC1..0xDF pseudo band is rewritten into real records before device fetch
• FP8 and dtype Encoding — the NEURON_ISA_TPB_DTYPE enum the record's dtype fields select from
• ISA Validator Architecture and Entry Tree — is_valid_enum, the OPCODE bitmasks, and the per-opcode legality predicates
• Per-Arch Instruction Validators (SUNDA / CAYMAN / MARIANA Deltas) — the arch-specialized validator copies and address-aperture deltas
• The Load Pipeline (Parse → Build → Stage → Relocate) — the stage step ib_read_write services; where record streams enter device HBM
• KBIN Structures — the KBIN container the record stream is carried in