PseudoCurProcessingRankID (0xDB)

NEURON_ISA_TPB_PSEUDO_CUR_PROCESSING_RANK_ID_STRUCT, opcode 0xDB (NEURON_ISA_TPB_OPCODE_PSEUDO_CUR_PROCESSING_RANK_ID). This is the device↔host NCCL rank bridge of the Cairo TPB ISA: a 64-byte pseudo-instruction the compiler emits at each step of a fine-grained collective so that, at runtime, the host collective layer can inject the per-iteration peer rank ID into a named scalar register on the engine. The compiled device program carries only opaque indices — which replica group, which ring, which iteration step, and a where-to-deposit register slot; it knows nothing about absolute rank numbers. The runtime ("RT/NCCL" in the header text) is what dereferences those indices against this process's actual rank and the NEFF's replica-group membership table, then materialises the resulting integer into the engine register file.

This page also documents its sibling GID_LOAD (0xDC) — NEURON_ISA_TPB_PSEUDO_GID_LOAD, "loads the current rank's GID into a register" — which shares the write-a-rank-identity-into-a-numbered-register idiom but delivers this rank's global ID rather than a peer's rank for an iteration.

All struct / enum / opcode facts below are read from the shipped clean C ISA headers of aws-neuronx-gpsimd-customop-lib 0.21.2.0 (arch neuron_cayman_arch_isa, banner ISA header for NC-v3 → CAYMAN = NC-v3) and a self-contained gcc probe that printed sizeof/offsetof for every field on all four arch variants. The headers, the device microcode binary libnrtucode_internal.so, the host loader libnrtucode.so, and instruction_mapping.json are all binary-derived artifacts of the package and are cited directly.

NOTE — pseudo-op, never executed as-is. 0xDB and 0xDC both live in the pseudo range 0xC1..0xDF. The OPCODE enum states the rule verbatim at aws_neuron_isa_tpb_common.h:263: "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." 0xDB = 0b110_11011 ✓, 0xDC = 0b110_11100 ✓. Confirming the rewrite is host-side: the device GPSIMD microcode (libnrtucode_internal.so) ships decode_* routines and SB2SB_Collective / POLL_SEM strings for the lowered legs, but has no cur_processing_rank or gid_load decoder string at all — these pseudo words are consumed and erased by the host before any NeuronCore sees them. (HIGH / OBSERVED.)

1 · PseudoCurProcessingRankID struct layout (byte-exact, compile-verified)

aws_neuron_isa_tpb_pseudo_cur_processing_rank_id.h:23-33. The struct is not __packed__ at the top level, so natural C alignment applies; total size is 64 bytes = one TPB instruction word. Six payload fields sit between the common header/events prefix and the tail padding.

(HIGH / OBSERVED — gcc probe over the real shipped header printed sizeof==64 and each offsetof above verbatim; reproduced for all four arches in §2.)

// gcc probe output, /tmp/gpsimd-p10/probe_real (real header, -I .../neuron_cayman_arch_isa/tpb)
// === rank_id (0xDB) ===
//   sizeof = 64
//   header 0  events 4  group_id 12  channel_id 14  iteration_id 16
//   dst_reg 18  reserved0 19  stream_id 20  reserved1 22
//   sub: HEADER=4  EVENTS=8  REG_NUM=1  OPCODE=1

CORRECTION — six payload fields, not three. A naive reading lists only group_id / iteration_id / dst_reg. The shipped header carries six payload fields in order — group_id(12), channel_id(14), iteration_id(16), dst_reg(18), reserved0(19), stream_id(20). channel_id (the ring index) and stream_id (the concurrent-stream selector) are present and real, and are exactly the vocabulary the v2 extension re-exposes (see TriggerCollective2 + Ext, which surfaces channel_id @ ext+4 and stream_id @ ext+6 as "NEW" knobs — they are new to the v2 trigger path, but the fine-grained rank-id op has always carried them). (HIGH / OBSERVED.)

Union binding. aws_neuron_isa_tpb_util.h:64 carries the member into the instruction union: NEURON_ISA_TPB_PSEUDO_CUR_PROCESSING_RANK_ID_STRUCT pseudo_cur_processing_rank_id;. instruction_mapping.json:105-107 binds the struct to its single opcode: "NEURON_ISA_TPB_PSEUDO_CUR_PROCESSING_RANK_ID_STRUCT" → ["NEURON_ISA_TPB_OPCODE_PSEUDO_CUR_PROCESSING_RANK_ID"]. The master ISA aggregate pulls the header in at aws_neuron_isa_tpb.h:41. (HIGH / OBSERVED.)

QUIRK — 0xDB has NO static validator. aws_neuron_isa_tpb_assert.h:12651 and aws_neuron_isa_tpb_assert_declaration.h:1119-1120 contain only the section marker //* pseudo_cur_processing_rank_id_assert.h with no body — there is no is_valid_pseudo_cur_processing_rank_id() predicate anywhere. In the per-opcode dispatcher debug_invalid_neuron_isa_instruction() (aws_neuron_isa_tpb_debug_assert.h:24), 0xDB falls into the shared case at line 442–443 whose body is check.result = true; // always pass pseudo ops, no checks in the ISA. The opcode otherwise appears only in classification sets — is_pseudo_op (instr_assert.h:594/621), neuron_isa_check_opcode_on_engine (instr_assert.h:675/856), and is_control_instruction (instr_assert.h:1143/1181) — never in a per-field check. Field validity (group_id range, dst_reg within the register count, iteration_id < num_ranks) is the runtime's responsibility, consistent with RT/NCCL filling the values. Contrast 0xDC, which does ship a real validator (§7). (HIGH / OBSERVED.)

2 · Layout is byte-identical on all four arch variants

The struct compiles to the exact same 64-byte layout on every arch the package ships — sunda (NC-v2), cayman (NC-v3), mariana (NC-v4), maverick (NC-v5). A diff across the four header copies shows the only differing line is the banner comment (ISA header for NC-v{2,3,4,5}); the typedef body is character-for- character identical. The same gcc probe compiled against each arch's include directory prints the identical offset set:

# probe_real built against each arch's tpb/ include dir:
sunda    : sizeof=64  group_id 12 channel_id 14 iteration_id 16 dst_reg 18 reserved0 19 stream_id 20
cayman   : sizeof=64  group_id 12 channel_id 14 iteration_id 16 dst_reg 18 reserved0 19 stream_id 20
mariana  : sizeof=64  group_id 12 channel_id 14 iteration_id 16 dst_reg 18 reserved0 19 stream_id 20
maverick : sizeof=64  group_id 12 channel_id 14 iteration_id 16 dst_reg 18 reserved0 19 stream_id 20

The opcode 0xdb is likewise stable on all four (common.h:289/287/295/301 for cayman/sunda/mariana/maverick respectively). (HIGH — v2/v3/v4 byte-grounded; v5/MAVERICK header-OBSERVED, struct interiors INFERRED wire-identical since the typedef bytes match — see the §2 register-range caveat for the one v5 knob that does differ.)

GOTCHA — dst_reg's valid range is NOT arch-invariant even though the layout is. dst_reg is a uint8_t, so the struct is identical across arches. But the architectural register file it indexes differs: NEURON_ISA_TPB_NUM_REGISTERS = 64U on sunda/cayman/mariana (common.h:31, NC-v2/3/4) versus 192U on maverick (NC-v5). A reimplementation that bounds-checks dst_reg must key the limit off the target arch, not off the struct. (HIGH / OBSERVED.)

3 · Base field types (aws_neuron_isa_tpb_common.h)

NEURON_ISA_TPB_HEADER (4 B, common.h:411-416) — shared by every TPB instruction, so 0xDB, 0xDC, 0xC8, 0xCB all open the same way:

typedef struct NEURON_ISA_TPB_HEADER {
    NEURON_ISA_TPB_OPCODE opcode;          // 1B, __packed__ enum  -> 0xDB / 0xDC
    uint8_t               inst_word_len;   // 1B
    uint8_t               debug_cmd;       // 1B
    uint8_t               debug_hint;      // 1B
} NEURON_ISA_TPB_HEADER;                   // sizeof == 4 (probe)

opcode is the __attribute__((__packed__)) enum NEURON_ISA_TPB_OPCODE (common.h:154-305), exactly 1 byte; the enumerator NEURON_ISA_TPB_OPCODE_PSEUDO_CUR_PROCESSING_RANK_ID = 0xdb is at common.h:289, flanked by …_PSEUDO_EXTENSION = 0xda (288) and …_PSEUDO_GID_LOAD = 0xdc (290).

NEURON_ISA_TPB_EVENTS (8 B, common.h:418-424):

typedef struct NEURON_ISA_TPB_EVENTS {
    NEURON_ISA_TPB_WAIT_MODE   wait_mode;       // 1B __packed__ (common.h:313)
    uint8_t                    wait_idx;        // 1B
    NEURON_ISA_TPB_UPDATE_MODE update_mode;     // 1B __packed__ (common.h:347)
    uint8_t                    update_idx;      // 1B
    uint32_t                   semaphore_value; // 4B
} NEURON_ISA_TPB_EVENTS;                        // sizeof == 8 (probe)

wait_mode is the full predicate enum (common.h:313-330): NONE=0x0, WAIT_FOR_SEM_{EQ,LT,LE,GT,GE}_IMM=0x1..0x5, WAIT_FOR_EVT_SET=0x6, WAIT_FOR_EVT_SET_THEN_CLEAR=0x7, the _REG variants 0x81..0x85, etc. update_mode (common.h:347-370) covers SEM_INC/DEC/ADD/SUB/WR on read- or complete-timing. For a fine-grained CCop these are how each lowered iteration rendezvouses with the producing/consuming engine — though, like the rest of the pseudo word, the runtime is free to rewrite them during lowering.

NEURON_ISA_TPB_REG_NUM (common.h:703):

typedef uint8_t  NEURON_ISA_TPB_REG_NUM;   // sizeof == 1 (probe) — names one engine register

CORRECTION — the engine register file is the 64-/192-entry GP file, not a 32-entry one. It is tempting to read dst_reg's target as the NEURON_ISA_TPB_MEM2D_DATA.registers[32] array (common.h:757), but that union is a 32-byte SIMD data-lane view (uint8[32]/uint16[16]/fp32[8] aliases), not the scalar register file. The scalar file is enumerated by NEURON_ISA_TPB_ENGINE_REGISTERS (common.h:…-1133) as R0..R63 plus the special engine registers PC_LO=128, PC_HI=129, DGE_TABLE_LO=130, DGE_TABLE_HI=131; the valid-write count is NUM_REGISTERS (64 / 192, §2). dst_reg names one of those GP registers. (HIGH / OBSERVED.)

4 · The dst_reg rank-injection mechanism

The header doc comment (aws_neuron_isa_tpb_pseudo_cur_processing_rank_id.h:14,20) gives the mechanism in two clauses:

"This instruction is used with CCop with Permute and PermuteReduce, it is lowered to TLD control instruction during runtime." "dst_reg: holds the current iteration processing rank ID, RT/NCCL feeds this value in during runtime."

Decoded as annotated pseudocode — note dst_reg is a static register number chosen by the compiler, never the rank value itself:

// FINE-GRAINED CCop RANK INJECTION  [register-injection model HIGH; exact lowered TLD encoding MED]
//
// COMPILE TIME (device program, emitted per unrolled iteration):
//   for (iter = 0; iter < num_ranks; ++iter) {
//       emit PseudoCurProcessingRankID{
//           .group_id     = G,        // which replica group  (uint16 @ off 12)
//           .channel_id   = ring,     // which CC ring 0/1     (uint16 @ off 14)
//           .iteration_id = iter,     // which step            (uint16 @ off 16)
//           .dst_reg      = Rk,       // a STATIC reg NUMBER   (uint8  @ off 18) -- "where to put the answer"
//           .stream_id    = S };      // which CC stream       (uint16 @ off 20)
//       ... per-iteration body that READS engine_reg[Rk] as the peer selector ...
//   }
//
// MODEL LOAD / RUNTIME (host RT + NCCL-equivalent collective layer):
//   my_rank   = this_process_rank_within(group_id /*G*/, channel_id /*ring*/);
//   peer_rank = collective_step_function(my_rank, iteration_id, ring, num_ranks);  // lives in NCFW/RT
//   //  the pseudo-op is REPLACED by a concrete TLD CONTROL instruction that
//   //  loads peer_rank into engine_reg[dst_reg]:
//   engine_reg[dst_reg] = peer_rank;     // materialised peer rank for THIS iteration
//
// DOWNSTREAM (same lowered sequence):
//   a DMA / SendRecv(0xCB) leg addresses engine_reg[dst_reg] as its peer selector.

Three structural facts ground the "lowered to TLD control instruction" phrase against the binary, rather than taking it on faith:

1. is_control_instruction(0xDB) == true (instr_assert.h:1143-1185, the 0xDB case at line 1181) — the ISA itself classes this pseudo-op as a control instruction, exactly the family the header says it lowers into. (0xDC is in the same set at line 1182.)
2. There is no per-field validator (§1 QUIRK), so the value placed in dst_reg is opaque to the static ISA — it is a host-supplied integer, not a statically-checkable encoding.
3. The device microcode does not decode 0xDB at all (intro NOTE), so the load it performs into engine_reg[dst_reg] must be emitted as one of the real control instructions the device does execute (a register move/load), i.e. the "TLD control instruction".

The per-iteration mapping (my_rank, iteration_id, channel_id) → peer_rank is the rank→peer stepping function, which lives in the NCFW ring/kangaring walk, not in these headers — see NCFW Main Dispatch Loop and Ring + Kangaring Collective. This pseudo-op is the compiler-visible interface to that firmware walk. (stepping-function residence HIGH; its exact arithmetic LOW — deferred to the NCFW pages.)

5 · The replica-group vocabulary (group_id / channel_id / stream_id)

These three fields are the shared collective-participant vocabulary; the identical glosses appear on TriggerCollective (0xC8) (group_id) and on the v2 extension (group_id/channel_id/stream_id). The trigger op names which collective runs on which group; this rank-id op walks the ranks of that group.

group_id (uint16, off 12) — a foreign key into host replica-group state. replica_groups is the XLA/JAX collective participant-set abstraction: a partition of all ranks into groups, where a collective with group_id = G runs independently within each group and the membership of G defines who exchanges data. The device pseudo-op carries only the index G; the membership table is host/NEFF data, injected by RT. The compiler-emitted NEFF subgraph def.json enumerates replica_groups alongside dma_queues and runtime_statebuffer_reservation, so the runtime dereferences group_id to a concrete rank list at load time. (HIGH — group_id gloss is verbatim header; replica_groups as host/NEFF metadata is OBSERVED in NEFF def.json.)

channel_id (uint16, off 14) — the ring index. Header verbatim: "if there are two CC topology rings, this value can only be 0 or 1." Dual-ring topologies split a collective's bandwidth across two simultaneous rings; channel_id picks which leg this iteration drives. The same ring concept reappears one layer down as the NCFW DMA-descriptor RING_ID phase — the ISA names the ring, NCFW sequences descriptors on it. (channel_id 0/1 gloss HIGH / verbatim; channel_id ↔ NCFW RING_ID identity MED — semantic match, not a literal cross-reference in headers.)

stream_id (uint16, off 20) — the concurrent-stream selector. Distinguishes independent in-flight collectives (CC streams). The host runtime enforces a one-to-one stream_id ↔ group_id mapping: the loader resolves it via enc_context::stream_map_get_stream_id(uint32_t) (libnrt @ 0xf6900), bounded by stream_id < NEC_MAX_STREAM_N, and rejects violations with the diagnostic "stream_id <-> group_id mapping violation. There should be a one-to-one …" (see TriggerCollective2 + Ext §host-resolution). (stream_id present/role HIGH; the precise libnrt resolver MED / CARRIED from the v2-ext lane.)

The collective-participant fields map cleanly onto the TriggerCollective (0xC8) and v2 vocabularies:

NOTE — fine-grained vs whole-collective. 0xC8 requests one whole collective over a group and lets the runtime pick ring/stream defaults; this rank-id op is the per-iteration primitive of the unrolled form, so it is the one op that carries iteration_id. The v2 path (0xD9/0xDA) is where the channel_id/ stream_id knobs become explicit on the trigger side too. (HIGH / OBSERVED.)

6 · CCop iteration over peer ranks, and which collectives use it

A fine-grained CCop (the header's term, from the Fine-Grained CCop Walrus implementation it references) does not issue one opaque collective — it is unrolled into per-rank iterations. For a group of N ranks the program holds a sequence in which iteration_id steps 0,1,…,N-1, and at each step PseudoCurProcessingRankID resolves the peer rank for that iteration into engine_reg[dst_reg], after which the body (a DMA / SendRecv to that peer, optionally a reduce) runs. This is the classic ring/scatter walk: each rank successively touches a different peer per step.

// FINE-GRAINED CCop LOOP (lowered shape)  [skeleton HIGH from header text; body composition MED]
//   for (iteration_id = 0; iteration_id < num_ranks; ++iteration_id) {
//       CurProcessingRankID(group_id, channel_id, iteration_id, dst_reg=Rk);  // peer -> reg[Rk]
//       SendRecv(0xCB){ .peer_id = reg[Rk], .is_send = ..., .tensor_id, .offset_bytes, .size_bytes };
//       // (optionally an in-DMA/CCE reduce for the *_REDUCE variants)
//   }

The header names the collective kinds it serves explicitly — "used with CCop with Permute and PermuteReduce". From NEURON_ISA_TPB_COLLECTIVE_TYPE (common.h:792-803):

Permutation collectives are inherently per-peer (each rank sends to one specific other rank), so they need an explicit "current peer rank" per iteration — which is precisely what dst_reg delivers. AllReduce / ReduceScatter / AllGather ring algorithms are likewise rank-stepped and naturally reuse the same primitive. (HIGH that the header names Permute/PermuteReduce; that 0xDB also backs the ring *_REDUCE/*_GATHER/*_SCATTER kinds = MED — INFERRED from the ring-walk structure, not stated for those ctypes in this header.)

The point-to-point counterpart that consumes reg[dst_reg] is PSEUDO_SEND_RECV (0xCB), aws_neuron_isa_tpb_pseudo_sendrecv.h, which carries uint32 tensor_id (off 12), uint64 peer_id (off 16), uint8 is_send (off 24), and uint64 offset_bytes/size_bytes (off 48/56) — sizeof==64, probe-verified. Its peer_id and the rank value the runtime materialises into dst_reg are the same NCCL rank vocabulary, so a fine-grained CCop composes as: for each iteration, CurProcessingRankID → reg[dst_reg] (the peer), then a SendRecv/DMA to that peer. (both structs HIGH / OBSERVED; that they compose into one lowered loop = MED / INFERRED from sibling ops.)

7 · GID_LOAD (0xDC) — "load current rank's GID into a register"

NEURON_ISA_TPB_PSEUDO_GID_LOAD_STRUCT, opcode 0xDC (NEURON_ISA_TPB_OPCODE_PSEUDO_GID_LOAD, common.h:290), aws_neuron_isa_tpb_pseudo_gid_load.h:32-41. Where 0xDB delivers a peer's rank for a collective iteration, 0xDC delivers this rank's own global ID (GID). It is the same "runtime writes a rank identity into a numbered scalar register" idiom, generalised to also allow an address destination.

// gcc probe output, /tmp/gpsimd-p10/probe_real
// === gid_load (0xDC) ===
//   sizeof = 64
//   header 0  events 4  dst_target 12  reserved0 13  dst 16  reserved1 24
//   dst union size = 8   GID_DST_REG size = 8   PSEUDO_ADDR8 size = 8   TARGET_TYPE size = 1
//   GID_DST_REG: reserved 0  regnum 7    -> regnum lands at off 16+7 = 23

Layout is byte-identical on all four arches and carries its own ISA_STATIC_ASSERT(sizeof == 64) at gid_load.h:41 (the rank-id struct has no such inline assert — a minor structural asymmetry). Bindings: util.h:65 pseudo_gid_load;, instruction_mapping.json:144, master include at aws_neuron_isa_tpb.h:54. (HIGH / OBSERVED.)

7.1 · The dst union

// gid_load.h:22-30  -- 8 bytes either way
typedef struct NEURON_ISA_TPB_PSEUDO_GID_DST_REG {
    uint8_t                reserved[7];   //  0 - 6   must be zero
    NEURON_ISA_TPB_REG_NUM regnum;        //  7       target scalar register (uint8)
} NEURON_ISA_TPB_PSEUDO_GID_DST_REG;

typedef union NEURON_ISA_TPB_PSEUDO_GID_DST {
    NEURON_ISA_TPB_PSEUDO_ADDR8       addr;        // 8B address form (common.h:616-623)
    NEURON_ISA_TPB_PSEUDO_GID_DST_REG reg_padded;  // 8B register form
} NEURON_ISA_TPB_PSEUDO_GID_DST;

When dst_target == REGISTER, dst.reg_padded.regnum (off 23) is the GP register the GID is written into, with reserved[0..6] (off 16..22) required zero. When dst_target == ADDRESS, dst.addr is a NEURON_ISA_TPB_PSEUDO_ADDR8 — the runtime's 8-byte tagged address (addr_reg / addr_var / addr_table / immediate variants discriminated by the high-byte marker, common.h:586-623) — so the GID can be written to SBUF / var memory instead of a register.

7.2 · 0xDC HAS a real static validator (asymmetry vs 0xDB)

Unlike 0xDB, GID_LOAD ships a compiled validator group ("Pseudo_Gid_Load_Block"), declared at assert_declaration.h:1207-1211 and defined with real bodies at aws_neuron_isa_tpb_assert.h:13044-13133. The header reproduces the same logic as a //-commented spec at gid_load.h:46-76; the implementation is in assert.h:

// is_valid_gid_load_dst  (assert.h:13065-13072) -- verbatim logic
bool is_valid_gid_load_dst(Inst i) {
    return (  ( i.pseudo_gid_load.dst_target == PSEUDO_GID_TARGET_TYPE_REGISTER
             && is_valid_register_write(i.pseudo_gid_load.dst.reg_padded.regnum)        // assert.h:753
             && is_zero_reserved_pseudo_gid_dst_reg(i.pseudo_gid_load.dst.reg_padded) ) // reserved[0..6]==0
           || ( i.pseudo_gid_load.dst_target == PSEUDO_GID_TARGET_TYPE_ADDRESS
             && ( gid_load_addr_reg(i.pseudo_gid_load.dst.addr)        // marker==ADDR_REG_BIT && reg_hi==0
               || gid_load_addr_var(i.pseudo_gid_load.dst.addr) ) ) ); // marker==VAR_ID_BIT
}

// is_valid_register_write  (assert.h:753-756) -- the register-range gate
bool is_valid_register_write(REG_NUM regnum) {
    return (regnum < NEURON_ISA_TPB_NUM_REGISTERS)              // < 64 (v2/3/4) or < 192 (v5)
        || regnum == NEURON_ISA_TPB_ENGINE_REGISTERS_DGE_TABLE_LO   // 130
        || regnum == NEURON_ISA_TPB_ENGINE_REGISTERS_DGE_TABLE_HI;  // 131
}

So 0xDC's regnum is validated to be a real GP register (< NUM_REGISTERS) or one of the two DGE-table engine registers, and the 7 reserved bytes must be zero. (HIGH / OBSERVED — full validator bodies present and compiled in assert.h.)

CORRECTION — 0xDC's validator is real, not merely commented. It is easy to conclude from the //-prefixed spec in gid_load.h:46-76 that the gid-load validator is only a comment. It is also a compiled C function at assert.h:13044-13133 (is_valid_pseudo_gid_load → is_valid_gid_load_dst → is_valid_register_write + is_zero_reserved_pseudo_gid_dst_reg). This is the sharpest contrast with 0xDB: the peer-rank op (0xDB) trusts the runtime entirely (no validator), whereas the own-GID op (0xDC) statically checks its destination. A reimplementation's verifier must apply the gid-load dst check but must not expect a cur-processing-rank-id check. (HIGH / OBSERVED.)

8 · Reimplementation checklist

To emit and lower 0xDB / 0xDC on a Vision-Q7-compatible engine:

1. Encode 0xDB as the 64-byte word of §1: header.opcode = 0xDB, inst_word_len per ISA, then events, group_id(12), channel_id(14), iteration_id(16), dst_reg(18), zero reserved0(19), stream_id(20), zero reserved1[42]. The 0xDB upper-3-bits 0b110 invariant must hold or NRT will not treat it as a pseudo-op.
2. Choose dst_reg as a static GP register number (< NUM_REGISTERS for the target arch: 64 for v2–v4, 192 for v5) — it is a slot, never the rank value.
3. At load time, for each iteration, resolve (my_rank, iteration_id, channel_id, group_id) → peer_rank (the NCFW/RT ring-step function) and replace the pseudo word with a control instruction that loads peer_rank into engine_reg[dst_reg].
4. Compose the body: have the per-iteration SendRecv(0xCB) / DMA leg read engine_reg[dst_reg] as its peer_id selector; add the in-DMA/CCE reduce for *_REDUCE kinds (see TriggerCollective §5).
5. Resolve group_id against the NEFF def.json replica_groups table to the concrete rank list; enforce the one-to-one stream_id ↔ group_id mapping.
6. For 0xDC, set dst_target (REGISTER/ADDRESS), fill dst.reg_padded.regnum (off 23) with reserved[0..6]=0 or dst.addr with a valid PSEUDO_ADDR8 (addr_reg marker 0x80 & reg_hi==0, or addr_var marker), and run the §7.2 validator before accepting it. The runtime writes this rank's GID into the chosen destination.
7. Do not implement a device decoder for either op — both are erased host-side; the device only ever sees the lowered control/DMA legs.

For the full host trace (NEFF → load-time rewrite → TOP_SP cc_op SPAD program → ring/mesh/hier step → CCE leg), follow A Collective, End to End, where this op appears at step "0xdb PSEUDO_CUR_PROCESSING_RANK_ID (per-rank index)".

See also

• TriggerCollective (0xC8) — the whole-collective trigger that shares the group_id/channel_id replica-group vocabulary and lowers to the same TOP_SP/DMA legs.
• TriggerCollective2 + Ext (0xD9/0xDA) — the v2 path that re-exposes channel_id/stream_id and the host stream_id ↔ group_id resolver.
• SendRecv (0xCB) — the point-to-point peer_id transfer the fine-grained loop drives with reg[dst_reg].
• Collective-Type + cc_op Enum Reference — the full COLLECTIVE_TYPE (Permute/PermuteReduce/AllReduce/…) catalog.
• The Unified Collective-Communication Architecture — where the device↔host rank/world model and replica-group state live.
• NCFW Main Dispatch Loop and Ring + Kangaring Collective — the firmware that owns the (my_rank, iteration_id, channel_id) → peer_rank stepping function.
• A Collective, End to End — the orientation trace tying the rank loop into the full lowering pipeline.