The Penguin BIR Instruction Set + BIR→ISA Map

Scope. The complete penguin BIR Inst* class roster recovered byte-exact from libBIR.so (110 distinct concrete classes), and for each GPSIMD-relevant class: its operands, engine binding, the SundaISel codegen rule that lowers it, the resulting SundaISAInst, and the numeric TPB opcode. This page sits one level above the emit_*→opcode table: both the NKI emit_ surface and the XLA-HLO frontend converge on the same BIR Inst* nodes, so the BIR→ISA map here is the superset of the emit table. Numeric opcodes are reconciled against the device opcode roster in opcode-catalog-ledger.

Audience. A reimplementer building a Vision-Q7-compatible engine who needs to know, for any IR node the compiler can produce, which device byte it ultimately decodes to — and, conversely, which device opcodes the compiler never emits (firmware-internal).

All facts below derive from static analysis of the shipped neuronx-cc 2.24.5133.0+58f8de22 wheel: the stripped-but-symbol-bearing libBIR.so dynamic symbol table, the Cython SundaISAInst / SundaISel extension modules (whose qualnames survive in .rodata), and the committed device opcode ledger. Confidence is tagged per claim as HIGH/MED/LOW × OBSERVED/INFERRED/CARRIED, where OBSERVED = read directly from a shipped artifact this pass, CARRIED = OBSERVED in a cited in-repo page and reused.

1. Where Inst* lives in the lowering pipeline

The compiler walks five hops from a frontend op to a NEFF instruction byte. The BIR Inst* node is the pivot — the single IR vocabulary both frontends target:

  nki.isa.<op>  /  XLA-HLO op
    └─▶ penguin BIR  Inst<Name>            (libBIR.so: IRBuilder / createFromJson builds it)
          └─▶ [SundaISel.transformT<Name>Operator → <Name>CodeGen]   (the rewrite rule)
                └─▶ SundaISAInst.<Name>Op   (the target ISA inst — 33 of them)
                      └─▶ [libwalrus CoreVNGenImpl get_bytes]
                            └─▶ TPB opcode byte  +  NEURON_ISA_TPB_*_STRUCT descriptor → NEFF

Each concrete Inst* ships a fixed method-set [HIGH/OBSERVED]. Every class in the roster carries the same mangled-symbol skeleton in libBIR.so's dynamic table — this is how the roster is enumerable from a stripped binary:

  _ZN3bir Inst<Name> C1/C2 ...        constructor(name, BasicBlock[, fields])
  _ZN3bir Inst<Name> 14createFromJson ...    NEFF JSON round-trip (the canonical marker)
  _ZN3bir Inst<Name> 16getDefaultEngineEv    engine binding   (107 of 110 carry it)
  _ZTVN3bir<len>Inst<Name>E                   vtable           (110 of 110 carry it)

The JSON op-tag that lands in the NEFF *.json instruction stream is the class basename minus the Inst prefix (InstTensorTensor → "TensorTensor", InstQuantizeMx → "QuantizeMx"). That tag is exactly what the firmware decode in opcode-catalog-ledger keys against once SundaISel has bound it to an opcode.

NOTE — the three classes without getDefaultEngine. createFromJson and the vtable are present for all 110, but getDefaultEngineEv is present for 107. The three pure-control terminators that carry no engine are the structural-control nodes (InstTerminator/InstExit/InstReturn family); they expand to SEQ/RT-band pseudos, never to a data-engine opcode. [MED/INFERRED]

2. The full 110-class BIR Inst* roster

Recovered byte-exact this pass by joining the createFromJson and _ZTV symbol sets in libBIR.so (both yield the identical 110-element set; see §8 self-verify). Grouped by lowering family. [HIGH/OBSERVED]

A. POOL/DVE elementwise + reduce compute (30) InstTensorTensor InstTensorScalar InstTensorScalarPtr InstTensorScalarCache InstTensorScalarAffineSelect InstTensorReduce InstSelect InstCopyPredicated InstReciprocal InstMemset InstIota InstDropout InstExponential InstPool InstMax InstMaxIndex InstMaxIndexAndMatchReplace InstMatchReplace InstStreamShuffle InstStreamTranspose InstRangeSelect InstNonzeroWithCount InstGetSequenceBounds InstDveReadAccumulator InstAbstractCopy InstTensorCopy InstTensorCopyDynamicSrc InstTensorCopyDynamicDst InstGenericRelu InstTongaReduceMacroSymbolic

B. Gather / indirect / generic load-store (9) InstGather InstIndirectCopy InstIndirectLoad InstIndirectSave InstIndirectSaveAccumulate InstGenericIndirectLoad InstGenericIndirectSave InstGenericCopy InstGeneric

C. PE / systolic (4) InstMatmult InstMatmultBase InstMatmultMx InstMatmultSparse

D. ACT / activation (4) InstActivation InstActivationReadAccumulator InstReadActivationAccumulator InstLoadActFuncSet

E. Batch-norm (5) InstBNStats InstBNStatsAggregate InstBNBackprop InstBNBackprop2 InstBNGradients

F. MX microscaling (1) InstQuantizeMx

G. RNG (7) InstRng InstRand InstRand2 InstRandGetState InstRandSetState InstGetRandState InstSetRandState

H. Collective / distributed (8) InstCollective InstCollectiveCompute InstCollectiveSend InstCollectiveRecv InstGPSIMDSB2SB InstGetGlobalRankId InstGetCurProcessingRankID InstReadVarAddr

I. DMA / descriptor (9) InstDMA InstDMACopy InstDMABlock InstDMATrigger InstDMADescriptor InstDMADescriptorCCE InstDMADescriptorCopy InstDMADescriptorReplicate InstDMADescriptorTranspose

J. Load / store (4) InstLoad InstSave InstTensorLoad InstTensorSave

K. Control / terminator / register (24) InstNoOp InstHalt InstDrain InstEventSemaphore InstAllEngineBarrier InstCoreBarrier InstGroupResetSemaphores InstTensorCompletion InstBranchHint InstCompareAndBranch InstUnconditionalBranch InstLoop InstDoWhile InstDynamicForLoop InstBreak InstExit InstReturn InstCall InstTerminator InstSwitchQueueInstance InstResetQueueInstance InstRegisterAlu InstRegisterMove InstDevicePrint

L. Kernel / custom-op containers (5) InstCustomOp InstInlineASMBytes InstBIRKernel InstNKIKernel InstNKIKLIRKernel

30+9+4+4+5+1+7+8+9+4+24+5 = 110.

CORRECTION (vs DX-CC-06 §9 reproduction recipe). The report's repro command nm libBIR.so | rg -o 'Inst…C2ERKNSt' | sort -u yields 109, not 110, because one class's constructor does not take a std::string reference as its first argument: InstDynamicForLoop (its ctor signature carries loop-bound operands first). Its createFromJson (_ZN3bir18InstDynamicForLoop14createFromJsonE…) and vtable (_ZTVN3bir18InstDynamicForLoopE) are both present, so it is the 110th concrete class. Ground the count on createFromJson or _ZTV, not on the C2ERKNSt constructor pattern — the latter silently drops InstDynamicForLoop. [HIGH/OBSERVED]

NOTE — delta vs DX-CC-01's "73". The "73" was the GPSIMD/compute-touching subset (families A–I minus the control spine K and the load/store J). The full 110 adds the 24-class control spine (K), the 4 load/store (J), and the 5 kernel containers (L). No contradiction — same hierarchy at two granularities. [HIGH/INFERRED]

GOTCHA — Instruction and InstructionBasicBlockHolder are not in the 110. A naive vtable grep _ZTVN3bir…Inst…E returns 112; two of those (_ZTVN3bir11InstructionE @ 0x8fced8, _ZTVN3bir27InstructionBasicBlockHolderE @ 0x8fbfc0) are the abstract base classes of the hierarchy, not concrete leaf instructions. Filter them out (rg -v '^Instruction') to land on the exact 110. [HIGH/OBSERVED]

3. The complete BIR Inst*→ISA→TPB-opcode map

Columns: BIR class [engine] | operands / semantics | SundaISel rule | SundaISAInst | TPB opcode (numeric) | conf. Opcode numerics are the uppercase-hex values from opcode-catalog-ledger; every numeric below was cross-checked against that ledger this pass (engine column too). Where the report and ledger differ, a CORRECTION row follows.

3.1 POOL/DVE elementwise + reduce

3.2 Gather / indirect / generic

3.3 PE / ACT / Batch-norm / MX / RNG

CORRECTION (MX forward direction, vs the earlier 0x7B reading; CARRIED from DX-CC-03). InstQuantizeMx is the forward pack and lowers to 0xE3 QUANTIZE_MX on DVE, not 0x7B. 0x7B TENSOR_DEQUANTIZE (POOL) is the inverse dequant and has no forward BIR producer — it is the firmware-internal dequant the MX matmul consumes (§4.3a). The ledger confirms 0xE3 QUANTIZE_MX | DVE and 0x7B TENSOR_DEQUANTIZE | POOL. [HIGH/OBSERVED]

NOTE — InstMatmultMx is the only PE-MX forward producer. The MX matmul itself consumes packed weights; the per-element dequant is firmware-internal. So the MX path has exactly two host-visible BIR producers: InstQuantizeMx (0xE3, the activation pack) and InstMatmultMx (0x09/0x0A or folded 0x01/0x02). [HIGH/INFERRED]

3.4 Collective / DMA / Load-store / Control / Containers

CORRECTION (engine of 0xF1, vs DX-CC-06 §I grouping). DX-CC-06 lists InstDMADescriptorTranspose's 0xF1 DMA_GATHER_TRANSPOSE under the DMA family, which reads as an NX/DMA opcode. The committed ledger assigns 0xF1 to the POOL engine (0xF1 | DMA_GATHER_TRANSPOSE | POOL). Treat 0xF1 as a POOL-engine gather-transpose, not an NX/DMA descriptor, for engine-fan-in accounting. [MED/OBSERVED]

3.5 The fan-in: one BIR class, many opcodes

InstTensorTensor is the canonical polymorphic node — a single class whose opcode is selected at SundaISel by {dtype, engine, gen}, the mechanism analysed in dtype-engine-fanin-synthesis:

The same one-to-many shape applies to InstCopyPredicated (0x72/0x99/0xEA/0xE8), InstActivation (0x21/0x25/0x26), and InstTensorReduce (0x42/0x52 plus transpose 0x83/0x84 and cumulative 0x4E/0x5E).

4. The two boundary sub-rosters

4.1 The 2:1 splits — one concept, two BIR classes, two opcodes [HIGH/OBSERVED]

These look like a single operation but are genuinely two distinct BIR classes lowering to two distinct ISA insts:

• InstGather (0x68 PoolGather) vs InstIndirectCopy (0xE7 IndirectCopy). Both .PoolGather and .IndirectCopy qualnames are present in SundaISAInst.so — two classes, two ISA insts, two opcodes. InstGather is within-partition u32 (OOB→0); InstIndirectCopy is 8-core/16-part u16 with a ≤4096 bound. Not a name trap — a real BIR-level split. The flag-routed variants (InstCopyPredicated → four device forms; InstMatmultMx → v4 pair vs v5 folded) are one class with an opcode picked by flag.

4.2 BIR ops with NO direct opcode — macro / fan-out [HIGH/OBSERVED]

A SundaISel transform expands one BIR node into many ISA insts:

• InstCollective (AllReduce/AllGather/ReduceScatter/AllToAll) → Tiled*Op ISA insts → PSEUDO 0xC7/0xC8/0xD9 collective triggers + DGE/CCE descriptors. Relocated to a TOP-SP instruction series at load. No single device opcode. See fused-cc-lowering for the CC expansion.
• InstCollectiveCompute → CCE_OP descriptor (the reduce leg) — no standalone opcode.
• InstTongaReduceMacroSymbolic → CrossLaneReduce 0x7C/0x7D + a loop (LegalizePartitionReduce). One symbolic BIR node, a reduce-loop micro-schedule.
• InstBIRKernel / InstNKIKernel / InstNKIKLIRKernel and the Tiled fused-CC ops (TiledRmsNormOp / TiledSoftmax(Dx/RSum)Op / TiledNativeKernel{Attention,MLP,QKV,RMSNormQuant}): each is one BIR kernel node that RmsNormCodeGen / SoftmaxCodeGen / codegenTiledCCOp* expands into a sequence of POOL/ACT/DVE opcodes. The kernel is a micro-schedule, not an opcode.
• InstGeneric* (generic load/store/copy): conditionally lower to 0x68 gather (when can_lower_generic_load_to_gather) OR 0xBB DMA_INDIRECT OR a copy — the opcode is a SundaISel decision, not a fixed bind.
• Structural control (InstLoop/InstDoWhile/InstCall/InstReturn): unrolled or expanded at tiling; the FUNCTION_BEGIN/CALL/RETURN pseudos (0xD1–0xD3) are RT-band, relocated. No data-engine opcode.

4.3 Device opcodes with NO BIR producer — firmware-internal [HIGH/OBSERVED + CARRIED]

Opcodes in the device roster that no BIR Inst* emits — materialized by firmware:

• (a) Inverse MX dequant 0x7B TENSOR_DEQUANTIZE (POOL). The MX matmul and the dequant kernel consume packed MX data; the compiler emits only the forward InstQuantizeMx → 0xE3. 0x7B has no forward BIR producer.
• (b) The cptc codec impls (inside 0xE4 CONV_LUT_LOAD). dtype-selected sub-paths inside the 0xE4 dispatcher — no opcode, no BIR class; InstLoadActFuncSet reaches the dispatcher entry, the impls are internal.
• (c) BN param-loads 0x64/0x66/0x8E. The 256-reciprocal RAM load (0x66 LOAD_PARAMETER_RAM) and param2 load are firmware-internal sub-steps of the BN codegen, not distinct BIR nodes (the BIR has only InstBNStats/Aggregate/Backprop/ Gradients).
• (d) 0x25 ACTIVATE2 / 0x26 ACTIVATE_MULTIPASS — reached as lowering targets of InstActivation by a dtype/multipass flag (a fan-in like §3.5), no distinct BIR class.
• (e) Maverick int-wide 0xF3/0xF4 + SUNDA bf16 0x8A–0x8F — pure dtype-routed targets of InstTensorTensor/InstTensorReduce. No named BIR class, no emit_ name.
• (f) The NX SEQ sync ops 0xB1–0xB4/0xB6 (SET_ORDERING_MODE, MOVE_SHAPE, POLL_SEM, TEST_EVENT_SEM, COMPACT_CONTROL_INST) — SEQ/NX-firmware control with no host op. (The controllable sync ops 0xA0/0xA1/0xA2/0xB0 do have BIR producers: InstEventSemaphore/InstHalt/InstDrain/InstGroupResetSemaphores.)
• (g) The dormant band (0x04 WEIGHT_MASK, 0x05 WEIGHT_SHIFT, 0x4A–0x4C REG_LOAD/STORE/SHUFFLE, 0x73 ROI_ALIGN, 0x81 JPEG_DECODE) — legacy tonga opcodes the current BIR never emits.

NOTE — the RT-pseudo band 0xC1–0xDF is the inverse case. These 31 PSEUDO opcodes do have BIR producers (InstDMATrigger→0xC1 DMATRIGGER, InstCollective→0xC7/0xC8/0xD9, InstReadVarAddr→0xC9, InstSwitch/ResetQueueInstance→0xCF, InstGetGlobalRankId→0xDC, InstGetCurProcessingRankID→0xDB, InstAllEngine/CoreBarrier→0xD5/0xD8, InstTensorCompletion→0xDE) but are relocated by libnrt to real WRITE/SEMAPHORE/descriptor-trigger instructions before any device decode. So: BIR producer YES, device-firmware decode NO — the opposite of the §4.3 set. [HIGH/CARRIED]

5. The SundaISel rewrite rule set

The BIR→ISA rewrite rules read byte-exact from SundaISel.so .rodata qualnames this pass: 34 transform* methods + 21 *Codegen/*CodeGen methods. The pattern is two-phase — transformT<X>Operator canonicalizes the BIR-DAG, <X>CodeGen emits the ISA inst. [HIGH/OBSERVED]

The four SundaISel predicate symbols (OBSERVED in SundaISel.so this pass) that drive the conditional routes:

• can_lower_generic_load_to_gather — turns InstGeneric(Load) → 0x68 gather when the index pattern is within-partition u32; otherwise falls to 0xBB DMA_INDIRECT.
• can_use_dge — selects swdge / hwdge / none for the DMA-descriptor path.
• should_shard_offloaded — whether an offloaded mem op is sharded across cores.
• replaceLastAPAddrWithIndex — the indirect-addressing rewrite (the gather/scatter access-pattern index splice).

NOTE — InferIntrinsicOnCC / LegalizePartitionReduce are CARRIED, not OBSERVED. These two transform names appear in the lowering narrative but did not surface as literal .rodata strings in SundaISel.so this pass — they likely live in a separate transforms/Lower*.so or are inlined. The opcode results they produce (0x21 relu fold; 0x7C/0x7D cross-lane reduce) are OBSERVED via the ledger. Treat the rule names as [MED/CARRIED] pending a Lower*.so sweep. [MED/INFERRED]

5.1 Annotated codegen — the gather lowering

transformTGatherOperator → GatherCodegen is the cleanest non-trivial rule and the one whose 2:1 split (§4.1) is most instructive. The reconstructed shape:

/* SundaISel: lower InstGather  ->  SundaISAInst.PoolGather  ->  TPB 0x68 GATHER
 * Symbols: SundaISel.transformTGatherOperator (.rodata qualname, this pass)
 *          SundaISel.GatherCodegen / SundaGatherCodegenT
 *          predicate SundaISel.replaceLastAPAddrWithIndex (the AP index splice)
 *          target   SundaISAInst.PoolGather (qualname OBSERVED in SundaISAInst.so)
 * Device:  0x68 GATHER, POOL engine, S4D4_GT struct, idx u8/16/32, OOB -> 0
 */
ISAInst *gather_codegen(InstGather *bir) {
    /* (1) index dtype gate — InstGather carries the u32 within-partition contract;
     *     InstIndirectCopy is a *separate* BIR class for the u16 / <=4096 case (§4.1). */
    assert(bir->index_dtype == U32 && bir->within_partition);

    /* (2) splice the index access-pattern onto the last source AP dimension —
     *     this is replaceLastAPAddrWithIndex: the gather reads src[idx[lane]]. */
    AccessPattern src_ap = bir->src_ap;
    replace_last_ap_addr_with_index(&src_ap, bir->index_operand);

    /* (3) emit the PoolGather ISA inst — single opcode, no fan-out. OOB indices
     *     resolve to 0 in the device datapath (the GATHER struct's OOB-zero contract). */
    return new_PoolGather(/*opcode=*/0x68, src_ap, bir->dst_ap, /*oob=*/ZERO);
}

Contrast codegenGenericLoad, which is conditional — it only reaches 0x68 when the predicate holds, otherwise it emits a DMA descriptor:

/* SundaISel.codegenGenericLoad / SundaGenericLoadCodegen */
ISAInst *generic_load_codegen(InstGeneric *bir) {
    if (can_lower_generic_load_to_gather(bir))   /* within-partition u32 idx */
        return new_PoolGather(0x68, ...);        /* same target as InstGather */
    else
        return new_DMADescriptor(0xBB /*DMA_INDIRECT*/, ...);  /* fall to DGE path */
}

6. Cross-check — BIR roster vs the device opcode ledger

[HIGH]

1. Every GPSIMD-relevant BIR Inst* maps to a ledger opcode (or a §4.2 macro of them). The §3 map binds each compute/gather/PE/ACT/BN/MX/RNG/collective/DMA class to an opcode present in opcode-catalog-ledger. Zero BIR Inst* produces an opcode the ledger lacks. Verified this pass by look-up of all ~80 cited opcodes against the ledger's uppercase-hex table.
2. The ledger's NONE (decode-gap) opcodes re-classified by BIR producer. Most maintained-NONE opcodes (0x48 Reciprocal, 0x49 Memset, 0x72 CopyPredicated, 0x6A StreamShuffle, 0x85/0x86 CustomOp) have a BIR producer — the gap is a firmware decode debt, not a compiler-coverage gap. The 8 dormant-NONE opcodes (§4.3g) are the only ones with neither a BIR producer nor a maintained decode.
3. The three-dispatch-surface consistency holds at the BIR layer. The BIR JSON op-tag (== Inst basename) → SundaISAInst → TPB opcode byte threads cleanly into the .bin TPB-seq slot, the SEQ ASCII dispatch (index = opcode − 0x41), and the POOL Q7 kernel_info_table. The POOL-compute BIR ops (InstTensorTensor 0x41, InstTensorCopy 0x46/0x47, InstIota 0x7E, InstTongaReduceMacroSymbolic 0x7C/0x7D, InstGetSequenceBounds 0xBE, InstNonzeroWithCount 0xF2, 0x7B inverse-dequant, InstCustomOp/InstInlineASMBytes 0xF0) are exactly the kernel_info_table opcode column. [HIGH/CARRIED]

7. The extended emit_*→opcode table, now as BIR Inst*→ISA

DX-CC-01's 62 emit_ rows are the NKI-exposed subset of the 110 BIR classes. The §3 map adds the BIR ops the NKI surface does not name (reached only via XLA-HLO, dtype-routing, or macro internals):

• The PE sparse path (InstMatmultSparse → 0x06/0x07) — no emit_.
• The generic load/store/scatter (InstGeneric*/InstIndirect* → 0x68/0xBB/0x79) — the XLA generic-memory path, no emit_.
• The BN gradient/backprop (InstBNBackprop/Gradients → 0x65/0x63/0x94) — reached via nxd, not a direct emit_.
• The control/terminator spine (family K, 24 classes → SEQ/RT opcodes) — IR plumbing.
• The RNG state mgmt (InstRng/InstRand*/InstGetRandState → 0x4D/0x76/0x77/0x78) — no emit_.
• The kernel containers (InstBIRKernel/InstNKIKernel → fused-CC schedules) — the nki kernel-decorator path, not a single emit_.

So the §3 table is the complete BIR→ISA map: emit_ (62) + the XLA/macro/control delta (~48) = the 110-class roster, all bound to ledger opcodes or §4.2 macros.

8. Adversarial self-verify — the 5 strongest claims

Each claim is re-derived from the binary, attacked, and resolved. [HIGH/OBSERVED]

(1) The roster is exactly 110 — not 73, not 109, not 112. Three independent symbol families in libBIR.so must agree. createFromJson (the canonical concrete-class marker): 110. _ZTVN3bir…E vtables, after filtering the two abstract bases Instruction/InstructionBasicBlockHolder: 110. The two sets are byte-identical (diff empty). Attack: does the count depend on the cp-tag? The cp310/cp311/cp312 wheels are byte-identical per the version anchor, so the dynamic table is the same. Verdict: 110, triple-grounded. The "73" was the compute subset; "112" counts the abstract bases; "109" is the C2ERKNSt recipe artifact (claim 2).

(2) The C2ERKNSt repro recipe undercounts by exactly one — InstDynamicForLoop. createFromJson − C2ERKNSt = {InstDynamicForLoop}. Attack: is it a real concrete class or a parse artifact? Both _ZN3bir18InstDynamicForLoop14createFromJsonE… (T, text section) and _ZTVN3bir18InstDynamicForLoopE @ 0x8fd700 (V, weak vtable) are present — it is concrete; its constructor simply does not take a std::string& first. Verdict: the count must rest on createFromJson/_ZTV, never on C2ERKNSt.

(3) SundaISAInst ships exactly 33 *Op qualnames. strings SundaISAInst.so | rg -o 'SundaISAInst\.[A-Z][A-Za-z0-9_]+' | sort -u | wc -l = 33, including the split-confirming pair .PoolGather and .IndirectCopy (the §4.1 2:1 split) plus the five TiledNativeKernel{,Attention,MLP,QKV,RMSNormQuant} containers. Attack: are these runtime-constructed (so undercounted)? They are Cython class qualnames baked into .rodata at build, not dynamically named. Verdict: 33 ISA insts.

(4) InstQuantizeMx forward = 0xE3 (DVE), and 0x7B has no forward BIR producer. The ledger row 0xE3 | QUANTIZE_MX | DVE and 0x7B | TENSOR_DEQUANTIZE | POOL are distinct opcodes on distinct engines. InstQuantizeMx is the only quantize BIR class in the 110; there is no InstDequantizeMx. Attack: could 0x7B be reached via a flag on InstQuantizeMx? No — they are opposite directions on opposite engines, and the dequant is consumed by the MX matmul internally. Verdict: forward 0xE3, inverse 0x7B firmware-internal; corrects any 0x7B-forward reading.

(5) The collective family is a macro — no single device opcode. InstCollective lowers via transformAllReduce/AllGather/ReduceScatterOp → Tiled*Op ISA insts → the RT-pseudo triggers 0xC7/0xC8/0xD9, which the ledger documents as PSEUDO_* in the 0xC1–0xDF band (NRT-lowered, not device-decoded). Attack: is there a hidden single "collective" opcode? The ledger's 0xBF SB2SB_COLLECTIVE is the point-to-point leg (InstGPSIMDSB2SB/InstCollectiveSend/Recv), not the group collective. Verdict: group collectives are macro fan-outs to pseudos + DGE/CCE descriptors; only the SB2SB p2p leg has a real opcode.

9. Version anchors

• neuronx-cc 2.24.5133.0+58f8de22 — libBIR.so (the 110-class roster + the createFromJson/_ZTV symbols), SundaISAInst.cpython-311-*.so (33 *Op), SundaISel.cpython-311-*.so (34 transform* + 21 *Codegen/*CodeGen) all read this pass. cp310/cp311/cp312 byte-identical.
• The 110-class BIR roster is gen-agnostic — one IR for all targets. The per-gen opcode split (sunda / cayman / mariana / maverick) is applied at SundaISel/walrus, not at the BIR layer; no silicon-generation fact is inferred from any compiler descriptor.

See also

• The GPSIMD-Relevant Compiler Map + emit_*→opcode — the NKI-exposed subset this page is the superset of.
• SundaISel Deep-Dive — the BIR-DAG match / canonicalize internals.
• Fused-CC / nkilib Kernel Lowering — how the §4.2 kernel macros expand to POOL/ACT/DVE schedules.
• Dtype/Engine/Gen Fan-In + CC-Lane Synthesis — the §3.5 one-class-many-opcode mechanism.
• Opcode Catalog Ledger — the device-side opcode roster every numeric here is reconciled against.