SENDRECV (point-to-point, 0xCB)

NEURON_ISA_TPB_PSEUDO_SENDRECV_STRUCT, opcode 0xCB (NEURON_ISA_TPB_OPCODE_PSEUDO_SEND_RECV). This is the point-to-point send+recv primitive of the Cairo TPB collective ISA: a single 64-byte pseudo-instruction that moves size_bytes of one local buffer to (or from) exactly one remote rank. It is the irreducible building block of every ring and mesh collective — a ring step is a SENDRECV to a neighbor — and the simplest member of the family: one directed State-Buffer→State-Buffer copy, no neighbor-chaining, no arithmetic fold.

The headline structural fact, and the one that catches every re-implementer off guard: there is no SEND opcode and no RECV opcode. There is one opcode, 0xCB, and one struct, carrying a single 1-byte is_send flag that selects the leg — is_send=1 ⇒ SEND (TX), is_send=0 ⇒ RECV (RX). The host runtime re-expands that one bit into the two distinct host op-kinds ENC_SEND=5 / ENC_RECV=6 and emits a single directed leg of the same SB2SB/RDMA data path documented in RDMA Cross-Die SBUF→SBUF P2P and SB2SB Remote-Copy Kernel (0xBF).

Provenance. Everything below is grounded in three shipped, redistributable artifacts: the clean C ISA headers (aws-neuronx-gpsimd-customop-lib 0.21.2.0, compile-verified struct + enum bodies), instruction_mapping.json (the struct→opcode binding), and libnrt.so (BuildID 8bb57aba0fb2e0035f1d88e9fc4fb3e7387c102e, x86-64, not stripped, with DWARF — the host runtime that lowers the pseudo-op). All addresses are file offsets; in this image .text, .rodata and .data all satisfy VMA==file-offset (no delta — readelf -SW: .text@0x3dbc0, .rodata@0x7cf000, .data@0xc07e00).

1. The opcode — one opcode, not SEND+RECV

From aws_neuron_isa_tpb_common.h (the NEURON_ISA_TPB_OPCODE enum) and instruction_mapping.json. [HIGH / OBSERVED]

// aws_neuron_isa_tpb_common.h:273  (cayman/NC-v3; identical across all 4 archs)
NEURON_ISA_TPB_OPCODE_PSEUDO_SEND_RECV  = 0xcb,

Neighbours in the pseudo-op opcode space (same header, surrounding lines):

The struct→opcode binding (verbatim, instruction_mapping.json, sub-map struct2opcode, present in all four arch copies cayman/mariana/maverick/sunda):

"NEURON_ISA_TPB_PSEUDO_SENDRECV_STRUCT": [
    "NEURON_ISA_TPB_OPCODE_PSEUDO_SEND_RECV"
]

NOTE. The binding lives under the struct2opcode key (not struct2pseudo_opcode) in every copy. A jq '.struct2pseudo_opcode.…' lookup returns null and would mislead you into thinking the mapping is absent.

Why 0xCB is not a HW instruction. The pseudo-op class rule is a verbatim comment in the same enum (common.h, at …PSEUDO_DMATRIGGER = 0xc1):

"all pseudo instructions have upper three bits of the opcode equal to 0b110. Pseudo instructions are generated by compiler and translated into non-pseudo HW instructions by NRT."

0xCB has top bits 0b110 (the 0xC0..0xDF window), so it is compiler-emitted and lowered host-side by NRT into the SB2SB (0xBF) + DMA + semaphore legs (§5). The NEURON_ISA_TPB_OPCODE enum is NEURON_ISA_PACKED (__attribute__((__packed__))), so the opcode occupies exactly 1 byte in the instruction header.

2. The struct — NEURON_ISA_TPB_PSEUDO_SENDRECV_STRUCT (64 B, byte-exact)

Verbatim from aws_neuron_isa_tpb_pseudo_sendrecv.h. The struct body is byte-identical across cayman (NC-v3) / mariana (NC-v4) / maverick (NC-v5) / sunda (NC-v2) — diff of tail -n +5 is empty for every pair; only the "ISA header for NC-v{2,3,4,5}" comment differs. [HIGH / OBSERVED + COMPILE-VERIFIED]

typedef struct NEURON_ISA_TPB_PSEUDO_SENDRECV_STRUCT {
    NEURON_ISA_TPB_HEADER header;        // 4    ( 0 -  3)
    NEURON_ISA_TPB_EVENTS events;        // 8    ( 4 - 11)
    uint32_t              tensor_id;     // 4    (12 - 15)
    uint64_t              peer_id;       // 8    (16 - 23)
    uint8_t               is_send;       // 1    (24)
    uint8_t               reserved0[23]; // 23   (25 - 47)
    uint64_t              offset_bytes;  // 8    (48 - 55)
    uint64_t              size_bytes;    // 8    (56 - 63)
} NEURON_ISA_TPB_PSEUDO_SENDRECV_STRUCT;

Compile-verified (gcc -Wpadded against the real shipped header, this session):

sizeof(SENDRECV)=64
header@0  events@4  tensor_id@12  peer_id@16  is_send@24  reserved0@25  offset_bytes@48  size_bytes@56
sizeof(HEADER)=4   sizeof(EVENTS)=8   sizeof(OPCODE)=1   -Wpadded clean → no implicit padding

The two sub-structs (common.h:411 / common.h:418, verbatim):

typedef struct NEURON_ISA_TPB_HEADER {      // 4 B
    NEURON_ISA_TPB_OPCODE opcode;           //   = 0xCB (1 B, packed enum)
    uint8_t               inst_word_len;    //   instruction word length
    uint8_t               debug_cmd;        //   debug command
    uint8_t               debug_hint;       //   debug hint
} NEURON_ISA_TPB_HEADER;

typedef struct NEURON_ISA_TPB_EVENTS {      // 8 B  — the inline completion semaphore
    NEURON_ISA_TPB_WAIT_MODE   wait_mode;       // 1  WAIT_MODE_{NONE,SEM_EQ/LT/LE/GT/GE_IMM}
    uint8_t                    wait_idx;        // 1  which of 256 HW semaphores to wait on
    NEURON_ISA_TPB_UPDATE_MODE update_mode;     // 1  how to update the completion sema
    uint8_t                    update_idx;      // 1  which HW semaphore to update
    uint32_t                   semaphore_value; // 4  the wait/update immediate
} NEURON_ISA_TPB_EVENTS;

2.1 Field semantics

Where is dtype? Where is "count"? There is no dtype field and no element-count field in the struct. The element dtype is carried by the resolved tensor descriptor and flows into the lowering as the SDMA_DTYPE argument of the enc_primitive constructor (the demangled ctor signature ends …, SDMA_DTYPE, int, int, … — OBSERVED, §5/§7). The count is expressed as size_bytes; element-count = size_bytes / dtype_size. Crucially the ctor carries SDMA_DTYPE (the copy dtype), not SDMA_CCETYPE (the reduce op) — the structural proof of no-reduce (§7).

GOTCHA. Do not look for a src/dst pair. The single tensor_id names the local buffer for both legs; the remote endpoint is named only by peer_id, and the remote address is computed at runtime (§4) — it is never carried in the instruction. A SEND reads tensor_id+offset_bytes locally and writes the peer's SBUF; a RECV writes tensor_id+offset_bytes locally from the peer's SBUF.

NOTE — why "one struct + is_send" instead of two opcodes. SEND and RECV are a symmetric P2P pair: both legs share an identical layout (peer / buffer / offset / size / semaphore) and differ only in transfer direction. One opcode plus one direction bit is the ISA-economical encoding; the host re-expands is_send into the two leg-specific kinds ENC_SEND/ENC_RECV for distinct lowering. (Encoding OBSERVED; the rationale is INFERRED — MED.)

3. The host op-kind — ENC_SEND=5 / ENC_RECV=6

Read directly from the libnrt.so DWARF this session (enum DIE <0x3bbd0>, name enc_op_type, DW_AT_encoding = unsigned, DW_AT_byte_size = 4; every enumerator's DW_AT_const_value re-read). [HIGH / OBSERVED] This is the same enum and the same values committed in TriggerCollective2 + Ext and Collective-Type + cc_op Enum Reference:

ENC_ALLGATHER               = 0      ENC_PERMUTE                 = 8
ENC_ALLREDUCE               = 1      ENC_PERMUTE_REDUCE          = 9
ENC_BROADCAST               = 2      ENC_PERMUTE_IMPLICIT        = 10
ENC_REDUCE                  = 3      ENC_PERMUTE_REDUCE_IMPLICIT = 11
ENC_REDUCE_SCATTER          = 4      ENC_ALLTOALL_V              = 12
ENC_SEND                    = 5  ←   ENC_OP_INVALID = 13 (== ENC_OP_N = 13, sentinel)
ENC_RECV                    = 6  ←
ENC_ALLTOALL                = 7

The on-instruction is_send byte maps to the host op-kind:

// the host re-encoding (HIGH for the values; the is_send→{5,6} write is host-side)
enc_op_type op = (sr->is_send != 0) ? ENC_SEND /*5*/ : ENC_RECV /*6*/;

The host validator that receives this op asserts it is exactly a SEND or a RECV. The .rodata assert string is verbatim at file-offset 0x7f2008:

send_recv_op_type == ENC_SEND || send_recv_op_type == ENC_RECV

referenced from enc_metaring_primitive::__compose_p2p_channel (the reject path of the (op-5)<=1 guard, §5). [HIGH / OBSERVED]

NOTE — SEND/RECV are not a COLLECTIVE_TYPE variant. Unlike the reducing collectives, SEND and RECV have no entry in the on-instruction ISA COLLECTIVE_TYPE enum (range 0x0..0x9; see collective-enums.md). The point-to-point primitive lives entirely in the dedicated 0xCB opcode plus the host enc_op_type — it is not a ctype of the 0xC8/0xD9 triggers.

4. The PEER → SoC-address computation

The instruction carries only a peer_id (a u64 rank). Turning that into a cross-die SoC byte address is a three-link chain; each link is sourced and confidence-tagged.

4a. Peer-rank resolution — get_peer_id

get_peer_id(enc_replica_group_info const&, int, int) @ 0x101500 (disasm OBSERVED, fn extent 0x101500..0x101674). For the 2-rank replica group that encodes a sendrecv pair, it reads the group's rank vector via end - begin:

0000000000101500 <get_peer_id>:
  101511:  mov  0x30(%rdi),%rbx      ; ranks.end
  101515:  sub  0x28(%rdi),%rbx      ; ranks.begin  → rbx = byte length of rank vector
  …                                  ; (4-byte ranks: a 2-element group ⇒ length 8)

The canonical P2P semantics: a 2-element replica group encodes the {sender, receiver} pair, and peer = the element that is not self. [HIGH / OBSERVED — the begin@0x28 / end@0x30 group-struct offsets are byte-exact.]

The 2-rank communicator is bootstrapped by enc_init_sendrecv_group @ 0x139870 (NRT_STATUS enc_init_sendrecv_group(encd_context*, const enc_replica_group_info*, enc_glb_comm*), fn extent …..0x139d80). It exchanges peer info over ncclBootstrapSend / ncclBootstrapRecv, then calls nccl_init_comm / enc_init_comm, and asserts g_device_id != peer_id (a rank cannot send/recv to itself; assert string @0x841e4f). The two leg labels it logs are verbatim:

0x841e66  send/recv group, tx to peer
0x841e9a  send/recv group, rx from peer

It also references init_sendrecv_intra_node_comm(int, int, enc_glb_comm*, enc_comm*) for the same-node case. [HIGH / OBSERVED]

4b. Peer rank → NeuronCore node — the per-arch encd helpers

libnrt.so carries a per-arch family of peer-node helpers (OBSERVED symbols, all four archs): encd_arch_get_d2d_neighbor_devid (cross-die D2D neighbor), encd_arch_get_rmtv_neighbor_devid (the remote-V peer), and the cayman_/mariana_/sunda_get_p2p_* set. The cayman one, cayman_get_p2p_pod_peer_node @ 0x25ca90, computes the peer node from a virtual-core id (disasm OBSERVED, byte-exact):

0000000000025ca90 <cayman_get_p2p_pod_peer_node>:
  25cab5:  mov  0x14(%rax),%esi     ; core node base
  25cab8:  shr  $1,%esi             ; node = core / 2   (the LNC2 core-pairing)
  25cabf:  cmp  $0x1,%r13d ; je … xor $0x1   ; ─┐ direction code (2nd arg r13d):
  25cac7:  xor  $0x3,%esi                     ;  │ 0 ⇒ same, 1 ⇒ ^1, 2 ⇒ ^2, 3 ⇒ ^3
  25caca:  xor  $0x2,%eax  ; cmp $0x2,%r13d   ; ─┘ — the neighbor selector
  25cad9:  call 25b690 <cayman_get_p2p_port_dst_node>  ; resolve dst node
  25cae3:  addl $0x8,(%rbx)         ; += 8 — offset into the pod

The direction code is the host analogue of the encd_neigh peer ordinal (LOCAL/NEXT/PREV/PEER_RMTV/…). [fn body + shift/xor math HIGH/OBSERVED; the exact direction-code ↔ encd_neigh ordinal binding MED — same vocabulary, ordinals not byte-equated.]

4c. NeuronCore → cross-die SoC address — the {CAYMAN_ID, DIE} bitfield

The peer's NeuronCore becomes the SB2SB descriptor's {routing_id / remote_routing_id, remote_tpb_idx}, which is folded into the high bits of the SoC address exactly as RDMA Cross-Die SBUF→SBUF P2P decodes it:

A remote peer sets {CAYMAN_ID, CAYMAN_ID_VALID=1, DIE} (plus the neighbor decoder's EXIT_DIE / NEIGHBOR_ROUTE) so the SB2SB copy traverses the io_d2d fabric to the peer's SBUF aperture; an intra-node peer keeps CAYMAN_ID_VALID=0. [the SoC bitfield HIGH/OBSERVED in rdma-cross-die; the routing_id→high-bits rewrite MED.]

Summary of the peer chain (per-link confidence):

peer_id (u64 rank)
  └─get_peer_id──────────────────────────────────────► peer rank             [HIGH/OBSERVED]
      └─encd_arch_get_{d2d,rmtv}_neighbor_devid /
        cayman_get_p2p_pod_peer_node─────────────────► peer NeuronCore node  [HIGH fn / map MED]
          └─SB2SB rdma_desc_gen {routing_id,
            remote_tpb_idx}────────────────────────► SoC dst_addr
                                                       {CAYMAN_ID,DIE,VALID}  [MED]
              └─SoC-address decode────────────────► remote SBUF byte addr    [HIGH/OBSERVED]

5. The lowering to SB2SB — SEND = TX leg, RECV = RX leg

5a. The composer — __compose_p2p_channel

enc_metaring_primitive::__compose_p2p_channel(int, node const&, node_next const&, enc_op_type) @ 0x175d20 (fn extent …..0x176646). Its entry proves the op must be SEND or RECV, then splits on direction (disasm OBSERVED, byte-exact):

0000000000175d20 <__compose_p2p_channel>:
  175d68:  lea  -0x5(%r8),%eax      ; op_type - 5
  175d6c:  mov  %r8d,%r14d          ; keep op_type in r14d
  175d6f:  cmp  $0x1,%eax
  175d72:  ja   1765fc <reject>     ;  ⇒ (op_type - 5) <= 1  ⟺  op ∈ {5,6} = {SEND,RECV}
  …                                 ;     reject path raises the §3 assert
  175d8c:  cmp  $0x5,%r14d          ; op == ENC_SEND (5) ?
  175d90:  je   1763c8 <SEND-leg>   ;  yes ⇒ SEND/TX setup; else fall through ⇒ RECV/RX

The two legs dispatch to two distinct enc_primitive methods (call sites are instruction-exact in the function body):

The channel-field strings the composer writes are OBSERVED in .rodata: channel->net_recv @0x842b8a, net_send @0x842b9c — the P2P channel composes a net_send or net_recv edge. [HIGH / OBSERVED — the (op-5)<=1 guard, the cmp $0x5 SEND split, and both direct_send / direct_recv call targets are instruction-exact.]

5b. The complete call set of __compose_p2p_channel (exhaustive)

The full, deduplicated callee set (from the IDA callees of 0x175d20, demangled). This list is the structural evidence and is reproduced complete:

enc_primitive::direct_send(enc_half_chunk_index)          ← the SEND/TX leg
enc_primitive::direct_recv(enc_half_chunk_index,b,b,b,b)  ← the RECV/RX leg
enc_primitive::advance_send(int) / advance_recv(int)
enc_primitive::post_recv(int)    / reset_send_credit()
enc_primitive::mark_step(bool)   / mark_complete() / mark_end()
enc_primitive::proxy_start_operation()
enc_metaring_primitive::get_ring_default_credits()
enc_primitive::enc_primitive(…, SDMA_DTYPE, int, int, …)  ← ctor carries the copy DTYPE
( + operator new/delete, memcpy, memmove, nlog_write, __assert_fail,
    std::__throw_bad_array_new_length )

These are the same SB2SB step primitives every ring/mesh leg uses (the enc_primitive family). On device they resolve to the SB2SB (0xBF) iDMA copy. [HIGH / OBSERVED]

5c. The device leg — the SB2SB TX/RX protocol

On device, direct_send / direct_recv become a single directed leg of the SB2SB (0xBF) handler decoded in SB2SB Remote-Copy Kernel and S3D3 Collective (0xBF). The leg runs, per role: pre-sync → program_window (map the remote SBUF into the local Q7 window) → rdma_desc_gen (build the iDMA BD ring including the LOCAL + REMOTE semaphore descriptors) → rdma_desc_start (the tail-pointer doorbell). The is_send field maps 1:1 onto the SB2SB role:

is_send=1 ⇒ ENC_SEND ⇒ direct_send ⇒ SB2SB TX leg
            left_pop WAITS for the RX signal, then advances its M2S tail (TDRTP_inc)
            by the descriptor count → launches the CME COPY that streams the SBUF bytes
            to the peer's SBUF over io_d2d.
is_send=0 ⇒ ENC_RECV ⇒ direct_recv ⇒ SB2SB RX leg
            advances its S2M tail (RDRTP_inc) to publish empty recv descriptors, then
            right_push SIGNALS the TX side to proceed.

This is the same TX=M2S / RX=S2M left_pop/right_push handshake documented in rdma-cross-die.md. [the device TX/RX protocol HIGH/OBSERVED there; the is_send→TX/RX-role binding HIGH via the §5a direct_send/direct_recv split.]

5d. The instruction handler — instr_col_translate_psr / instr_col_setup_psr

instr_col_translate_psr @ 0x2705f0 is the 0xCB translation entry; it runs the pipeline build_enc_ctx → enc_stream_map_set_stream_id → instr_col_setup_psr → instr_col_add (OBSERVED callees). instr_col_setup_psr @ 0x26df50 resolves the buffer and validates the peer (OBSERVED callees): lookup_memref_by_idx, mem_ref_to_addr, mem_ref_get_va, get_enc_tensor_from_collectives_tensor (tensor_id → memory-region → SoC VA) and encd_get_global_comm (the comm/peer context). It references all three guard strings directly (§8). [HIGH / OBSERVED]

6. The ring-step building block — a ring step is a sendrecv

enc_metaring_primitive::__compose_channel_ring(int) @ 0x177340 calls __compose_p2p_channel @ 0x175d20 (OBSERVED call edge — the call site is 0x177c50; the callers of 0x175d20 are exactly two: __compose_channel_ring and __compose_permute_channel). The ring algorithm builds each step's neighbor leg by composing the same P2P channel the standalone SENDRECV uses — the only difference is which neighbor (next_neigh / prev_neigh) the channel targets and whether the leg is a SEND or a RECV.

The failure string names a ring step literally a "send/recv" (verbatim @0x7f23d0):

[nec_dev %d] failed to compose ring send/recv

A ring all-reduce / all-gather is therefore a sequence of these neighbor send/recv legs, chained by the recv_sema ≥ target counted barrier (see Ring + Kangaring Collective and the Ring-Protocol Config Command); a mesh/kangaring step fans out the same primitive to up to three peer semaphores. [the ring=sequence-of-sendrecvs model HIGH from the call edge + the string; the exact per-(world-size, topology) leg schedule lives in the LX/NCFW Xtensa core — MED/LOW.]

NOTE — the one divergence between a standalone SENDRECV and a ring step. A ring step's RECV leg may, for the reducing collectives, be a recv_reduce_* (fold-on-receive) rather than a plain direct_recv. The standalone SENDRECV's RECV is always a plain direct_recv (§7). That is the only behavioural difference; the composition path is identical.

7. No-reduce — SENDRECV is a pure MOVE

Three independent OBSERVED facts, mutually reinforcing:

1. Exhaustive call set (§5b). __compose_p2p_channel's complete callee list contains no recv_reduce_*, no direct_reduce_*, and no SDMA_CCETYPE — only direct_send / direct_recv plus the credit / advance / mark bookkeeping. A reducing leg would call recv_reduce_copy / recv_reduce_send / __recv_reduce_write / direct_reduce_*; the P2P composer calls none of them. [HIGH / OBSERVED]

2. Leg-primitive signatures. The leg methods the composer calls take no SDMA_CCETYPE argument: direct_send(enc_half_chunk_index) and direct_recv(enc_half_chunk_index, bool,bool,bool,bool). The reduce primitives are a distinct symbol set that do take (SDMA_CCETYPE, reduction_type_t, …) — verified directly in the binary, e.g.:

   enc_primitive::recv_reduce_copy (enc_half_chunk_index, SDMA_CCETYPE, reduction_type_t, …)  @0x16b030
   enc_primitive::recv_reduce_send (enc_half_chunk_index, SDMA_CCETYPE, reduction_type_t, …)  @0x16ad70
   enc_primitive::__recv_reduce_write(enc_half_chunk_index, SDMA_CCETYPE, reduction_type_t, …)@0x16a0a0
   

   None of these appears in the P2P composer. [HIGH / OBSERVED — symbol signatures.]

3. The struct (§2) and the ctor. The struct has no reduce-op field, and the enc_primitive ctor in this path carries an SDMA_DTYPE (the element dtype, for the copy) — not an SDMA_CCETYPE (the reduce op). [HIGH / OBSERVED]

So SENDRECV performs a pure data copy of size_bytes from tensor_id@offset_bytes to/from peer_id; it never folds. This matches NCCL P2P semantics (ncclSend/ncclRecv are non-reducing). [HIGH / OBSERVED]

8. Error / completion / timeout handling

8a. Completion — the semaphore / doorbell model

There is no timeout field in the struct; completion is the semaphore-poll model. The inline events field (off 4) is the per-op semaphore: the op waits its inbound semaphore before issue and updates the completion semaphore when done. The handler enforces the update — verbatim .rodata @0x80f938, referenced from instr_col_setup_psr:

Collective instruction must increment semaphore when done

On device the SB2SB CME COPY's last BD increments the local dma_compl_sema and, via the REMOTE semaphore descriptor routed by routing_id across the die fabric, the peer's recv_sema. The increment target is the EVT_SEM engine's tpb_semaphores_inc window at +0x1800 (write-only atomic +=; 256 hardware semaphores, 4-B stride) — the same two-semaphore (local_sem / remote_sem) completion model decoded in rdma-cross-die.md. The next step waits its recv_sema ≥ target. [HIGH / OBSERVED — device side; "no timeout field" HIGH from the struct.]

8b. Errors (all strings verbatim, OBSERVED)

8c. NOTE — a second, host-side P2P path exists (do not conflate)

libnrt.so also exposes an "async sendrecv" host API — 11 exported nrt_async_sendrecv_* symbols (init, connect, accept, send_tensor, recv_tensor, test_request, flush, close, …) — that runs over libfabric/OFI (the aws-ofi-nccl plugin), gated by the env var NEURON_RT_ASYNC_SENDRECV_EXPERIMENTAL_ENABLED and a bootstrap port (both OBSERVED in .rodata). That is a host inter-node network transport, distinct from the on-device 0xCB pseudo-op decoded here (which lowers to the on-chip SB2SB iDMA). Do not implement the 0xCB device leg against the async API's namespace. [HIGH / OBSERVED — both exist; the distinction is OBSERVED via the separate symbol/string namespaces.]

9. Cross-arch stability

• The SENDRECV struct body is byte-identical across cayman (NC-v3) / mariana (NC-v4) / maverick (NC-v5) / sunda (NC-v2) — only the "ISA header for NC-v{2,3,4,5}" comment differs. [HIGH / OBSERVED]
• The opcode PSEUDO_SEND_RECV = 0xCB and the struct2opcode binding are present in all four instruction_mapping.json copies. [HIGH / OBSERVED]
• libnrt.so is a single binary (one runtime serves all gens); enc_op_type and the __compose_p2p_channel lowering are gen-independent. Only the PEER→node helpers are per-arch (cayman_/mariana_/sunda_* + the encd_arch_* family). [HIGH / OBSERVED]

PROVENANCE OF v5/MAVERICK. The maverick (NC-v5) header was byte-diffed and is identical to the others for this struct; the per-arch maverick_* peer-node helper interiors are header/symbol-OBSERVED only — flag any maverick-specific addressing math as INFERRED until byte-grounded.

See also

• RDMA Cross-Die SBUF→SBUF P2P — the device data path SENDRECV lowers to: routing_id→SoC fold, TX=M2S/RX=S2M, the +0x1800 two-semaphore completion.
• SB2SB Remote-Copy Kernel (0xBF) and S3D3 Collective (0xBF) — the on-device 0xBF iDMA byte-mover that each direct_send/direct_recv leg becomes.
• TriggerCollective2 + Ext (0xD9/0xDA) — the source of the shared enc_op_type (ENC_SEND=5, ENC_RECV=6, ENC_ALLTOALL_V=12) catalog.
• Collective-Type + cc_op Enum Reference — the COLLECTIVE_TYPE (0x0..0x9) catalog that, deliberately, has no SEND/RECV entry.
• ALL_REDUCE — the reducing collective (trigger opcode 0xC7) whose ring steps add a recv_reduce_* fold on top of the plain SENDRECV legs documented here.
• Ring + Kangaring Collective and Ring-Protocol Config Command — the firmware that sequences SENDRECV legs into a ring/mesh schedule.
• PseudoCurProcessingRankID (0xDB) — the per-step peer_id injection the fine-grained collective loop drives.
• A Collective, End to End — the orientation trace tying the pseudo-op into the full NEFF → host → device lowering pipeline.