TriggerCollective (0xC8)

NEURON_ISA_TPB_PSEUDO_TRIGGER_COLLECTIVE_STRUCT, opcode 0xC8 (NEURON_ISA_TPB_OPCODE_PSEUDO_TRIGGER_COLLECTIVE). This is the canonical collective-communication pseudo-op of the Cairo TPB ISA: a single 64-byte instruction word the compiler emits to request one whole collective (AllReduce / ReduceScatter / AllGather / All-to-All / Permute) over a replica group, which the host runtime then rewrites — before it ever reaches a NeuronCore — into a sequence of TOP_SP (engine_idx 5) control + DMA instructions. It opens the collective pseudo-op catalog because every other collective carrier (0xC7 AllReduce, 0xD9/0xDA v2, 0xBF SB2SB) is best understood as a specialization or a lowering target of this op.

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

NOTE — pseudo-op, never executed as-is. 0xC8 lives in the pseudo range 0xC0..0xDF. The OPCODE enum documents 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." 0xC8 = 0b110_01000 ✓. The device-side GPSIMD microcode (libnrtucode_internal.so) has no decoder for 0xC8 — it decodes the lowered legs (SB2SB_Collective 0xBF, POLL_SEM 0xB3), confirming the rewrite happens host-side. (HIGH / OBSERVED.)

1 · Struct layout (byte-exact, compile-verified)

aws_neuron_isa_tpb_pseudo_trigger_collective.h:21-35. The struct is not __packed__ at the top level, so natural C alignment applies; total size is 64 bytes = one TPB instruction word. The reserved1[15] gap is deliberate — the header comment at line 32 states it pushes the two trailing uint64_t offset fields to an 8-byte boundary: "place this here cause C will force byte alignment padding for 8 byte field".

(HIGH / OBSERVED — gcc probe over the single header printed sizeof==64 and each offsetof above verbatim.)

// gcc probe output, /tmp/gpsimd-p10/probe (single header, -I tpb/)
// sizeof STRUCT = 64
//   header 0  events 4  op 12  dtype 13  group_id 14
//   input_tensor_id 16  output_tensor_id 20  num_elements 24
//   ctype 32  reserved1 33  src_offset_elems 48  dst_offset_elems 56
//   sizeof HEADER=4 EVENTS=8 ALU_OP=1 DTYPE=1 CTYPE=1 OPCODE=1 ENGINE=1
//   OPCODE_PSEUDO_TRIGGER_COLLECTIVE = 0xC8   ENGINE_TOP_SP = 5

CORRECTION — field order vs the task brief. The brief listed group_id before op and implied ctype near the front. The shipped header order is op, dtype, group_id (offsets 12/13/14), and ctype sits after num_elements at offset 32, not at the front. num_elements is uint64_t (not uint16/uint32), and src/dst_offset_elems are present at offsets 48/56. The table above is the authoritative layout.

Union binding. aws_neuron_isa_tpb_util.h:48 carries the member into the instruction union: NEURON_ISA_TPB_PSEUDO_TRIGGER_COLLECTIVE_STRUCT pseudo_trigger_collective;. instruction_mapping.json:190-191 binds the struct to its single opcode: "NEURON_ISA_TPB_PSEUDO_TRIGGER_COLLECTIVE_STRUCT" → ["NEURON_ISA_TPB_OPCODE_PSEUDO_TRIGGER_COLLECTIVE"]. (HIGH / OBSERVED.)

QUIRK — no static validator. aws_neuron_isa_tpb_assert.h:13574 contains only the section marker //* pseudo_trigger_collective_assert.h with no body — there is no is_valid_pseudo_trigger_collective() predicate. Contrast the real HW op s3d3_collective (assert.h:19215+), which ships a full is_valid_s3d3_collective() block checking lnc_size_fmt, num_active_channels, tensor patterns, and reserved-zero. TriggerCollective is validated and lowered by the runtime, not by the static ISA assert layer — consistent with it being erased before hardware sees it. (HIGH / OBSERVED.)

2 · Base field types (aws_neuron_isa_tpb_common.h)

NEURON_ISA_TPB_HEADER (4 B, common.h:411-416):

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

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__
    uint8_t                    update_idx;      // 1B
    uint32_t                   semaphore_value; // 4B
} NEURON_ISA_TPB_EVENTS;                        // sizeof == 8 (probe)

wait_mode is a full semaphore/event predicate enum (common.h:313+): NONE=0x0, WAIT_FOR_SEM_EQ_IMM=0x1, ..._LT/LE/GT/GE_IMM=0x2..0x5, WAIT_FOR_EVT_SET=0x6, WAIT_FOR_EVT_SET_THEN_CLEAR=0x7. For a collective the events block is how the lowered TOP_SP program rendezvouses with the producing/consuming engines — but for the v1 pseudo-op the runtime overwrites most of it during lowering. (HIGH / OBSERVED.)

All field enums used by this struct are __attribute__((__packed__)), probe-verified to 1 byte each: ALU_OP=1, DTYPE=1, COLLECTIVE_TYPE=1, OPCODE=1, NEURON_ENGINE=1. (HIGH / OBSERVED.)

3 · Enum values

See Collective-Type + cc_op Enum Reference for the full catalog; the four enums this struct directly references are reproduced here.

NEURON_ISA_TPB_COLLECTIVE_TYPE — ctype @ off 32 (common.h:792-803):

NEURON_ISA_TPB_ALU_OP — op @ off 12 (common.h:939-1000). The float/generic encodings are the low byte; the integer-engine variants share the bit[7:6]==0b11 block. For a collective reduce the practically meaningful ops:

(op encoding HIGH / OBSERVED. Which subset is legal for a collective is MED / INFERRED — no static assert constrains it in these headers; the practical answer comes from the DMA reduce ops in §5.)

NEURON_ISA_TPB_DTYPE — dtype @ off 13 (common.h:722-739): INVALID=0x0, UINT64=0x1, INT8=0x2, UINT8=0x3, INT16=0x4, UINT16=0x5, BFLOAT16=0x6, FP16=0x7, INT32=0x8, UINT32=0x9, FP32=0xA, FP32R=0xB, INT64=0xC, FP8_EXP3=0xD, FP8_EXP4=0xE, FP8_EXP5=0xF. (HIGH / OBSERVED.)

NEURON_ISA_TPB_NEURON_ENGINE — context for "TOP SP" (common.h:139-146): PE=0, ACT=1, POOL=2, DVE=3, TPB_SP=4, TOP_SP=5. The collective lowers onto engine 5. (HIGH / OBSERVED — probe ENGINE_TOP_SP == 5.)

4 · group_id and the replica-group vocabulary

group_id (off 14, uint16_t) selects which set of ranks participate in the collective. The vocabulary is shared with the fine-grained CCop pseudo-op PSEUDO_CUR_PROCESSING_RANK_ID (0xDB, aws_neuron_isa_tpb_pseudo_cur_processing_rank_id.h), whose doc comment names the full set the runtime collective framework consumes:

NOTE — what v1 omits. The v1 TriggerCollective carries only group_id. It has no channel_id (ring index) and no stream_id — the runtime supplies topology defaults during lowering. The richer knobs (channel_id, stream_id, cc_buf_*, DMA priority_class, A2Av bitmaps) appear only on the v2 path (§6, TriggerCollective2 + Ext). (HIGH / OBSERVED — cur_processing_rank_id struct: group_id/channel_id/iteration_id/ dst_reg/stream_id at offsets 12/14/16/18/20.)

The runtime layer that fills per-rank / per-iteration values and replaces the pseudo-op with concrete TOP_SP/DMA instructions is the collective framework ("RT/NCCL" in the header text; on the management side the Neuron Collective FW). The host load-time rewrite that actually relocates the pseudo word and composes the TOP_SP cc_op SPAD program is traced end-to-end in A Collective, End to End. (group_id meaning HIGH; the specific framework name is MED / CARRIED from the NCFW lane, not these headers.)

5 · TOP_SP lowering — what 0xC8 becomes

The header doc comment (aws_neuron_isa_tpb_pseudo_trigger_collective.h:14-17) states the model verbatim:

"TriggerCollective expands to a series of instructions executed on TOP SPs, which uses DMA to perform an the operation. TriggerCollective is a pseudo-instruction generated by compiler. It is replaced by TOP SP instructions by the runtime."

The lowering chain, pieced together from the sibling headers:

// LOWERING MODEL  [chain to DmaTrigger+Write = HIGH; exact 1:1 expansion = MED]
//
//   TriggerCollective(0xC8, ctype, op, dtype, group_id, ...)        // compiler
//        |  host runtime load-time rewrite (NRT/KRT)
//        v
//   series of TOP_SP (engine 5) instructions, per replica-group iteration:
//        PollSem    (0xB3)  -- min/decrement neighbor semaphores  (§5.2)
//        DmaTrigger (0xC1)  -- writes the DMA queue tail pointer   (§5.1)
//        ... barrier / rank-id glue (0xD5/0xD8/0xDB) ...
//        v
//   DmaTrigger -> a Write to the DMA queue tail-pointer CSR (KRT)
//        v
//   DMA engine performs the inter-rank transfer; the reduce is done IN-DMA (§5.3)

5.1 · The data-movement primitive — DmaTrigger (0xC1)

aws_neuron_isa_tpb_pseudo_dma_trigger.h defines the lowering's data primitive. Its doc comment: "DmaTrigger initiates a DMA data-transfer, by writing to the DMA tail pointer. DmaTrigger is a pseudo-instruction generated by compiler. It is replaced by Write, by KRT." Layout (64 B):

typedef struct NEURON_ISA_TPB_PSEUDO_DMA_TRIGGER_STRUCT {
    NEURON_ISA_TPB_HEADER  header;            // 0  - 3   (opcode 0xC1)
    NEURON_ISA_TPB_EVENTS  events;            // 4  - 11
    int8_t                 dma_queue_name[32];// 12 - 43  ASCII queue name -> TOP_SP-visible DMA ring/CSR
    uint32_t               use_raw_count;     // 44 - 47  block_id is a raw descriptor count when set
    union { uint32_t block_id; uint32_t count; } count_union; // 48 - 51
    uint8_t                reserved0[12];     // 52 - 63
} NEURON_ISA_TPB_PSEUDO_DMA_TRIGGER_STRUCT;

So the DMA leg is kicked by writing the queue's tail-pointer CSR. The descriptor ring is named by the 32-byte ASCII dma_queue_name, which the runtime maps to a concrete TOP_SP-visible DMA queue. When use_raw_count is set, count_union.count is a literal descriptor count; otherwise block_id references a pre-built block of descriptors from the k-elf (plus an optional event). (mechanism HIGH; exact CSR register numbers are NOT in these headers — LOW.)

5.2 · The per-iteration sync accelerator — PollSem (0xB3)

aws_neuron_isa_tpb_ctrl_poll_sem.h defines PollSem (NEURON_ISA_TPB_OPCODE_POLL_SEM = 0xB3, common.h:255), described as "a specialized semi-custom operation to help Neuron accelerate polling semaphores on TOP_SP", explicitly tagged to the TRN1 48MB AllReduce performance optimization. This is the per-iteration rendezvous primitive the lowered collective loop runs. Layout (64 B):

typedef struct NEURON_ISA_TPB_CTRL_POLL_SEM_STRUCT {
    NEURON_ISA_TPB_HEADER      header;             // 0  - 3   (opcode 0xB3)
    NEURON_ISA_TPB_EVENTS      events;             // 4  - 11
    uint8_t                    n_read;             // 12       1..16 consecutive uint32 reads
    NEURON_ISA_TPB_REG_NUM     result;             // 13       GPR receiving the folded result
    NEURON_ISA_TPB_ALU_OP      op;                 // 14       fold op (currently REQUIRED == Min)
    uint8_t                    reserved1[1];        // 15
    NEURON_ISA_TPB_NEURON_ADDR addr;               // 16 - 23  uint64 NeuronAddr, must be %4==0
    NEURON_ISA_TPB_NEURON_ADDR res_writeback_addr; // 24 - 31  uint64 NeuronAddr, must be %4==0
    uint8_t                    reserved2[32];       // 32 - 63
} NEURON_ISA_TPB_CTRL_POLL_SEM_STRUCT;             // ISA_STATIC_ASSERT sizeof==64

Behavioral contract (header comment, paraphrased to pseudocode):

// PollSem semantics  [HIGH / OBSERVED — verbatim header doc + commented validator]
//   read n_read consecutive uint32 from addr (n_read <= 16)
//   acc = vals[0]; for i in 1..n_read: acc = AluOp(op, acc, vals[i])   // op == Min
//   if acc != 0:
//       for i in 0..n_read: *(res_writeback_addr + 4*i) = acc          // write n_read copies
//   GPR[result] = acc                                                  // TPB program branches on it
// e.g. n_read=6, op=Min, *addr=[8,3,5,4,4,10] -> result=3, writeback=[3,3,3,3,3,3]

The static validator (commented in the header) requires n_read in 1..16, op == AluOp::Min, both addresses 4-byte aligned, and all reserved bytes zero — i.e. "select the minimum of a number of semaphores and decrement all of them by that amount." This is exactly the AllReduce credit/scoreboard sync done once per ring step. (HIGH / OBSERVED.)

5.3 · The reduce is done in-DMA (CCE / DGE compute-op)

The op field is an ALU_OP, but the element-level reduction is realized by the DMA engine's compute-on-transfer capability, not a separate ALU pass:

// NEURON_ISA_TPB_DGE_COMPUTE_OP (common.h:837-843)  -- reduction-add DMA mode
//   NONE=0x00 (B=A, plain memcpy)  ADD=0x01 (B+=A)  MULTIPLY=0x02 (B*=A)
//   MAX=0x03 (B=max(A,B))          MIN=0x04 (B=min(A,B))
//
// NEURON_ISA_TPB_CCE_OP (common.h:1002-1006)  -- "Same encoding as SDMA CCE op"
//   ADD=0x00   (Multiply=0x01, commented)   MAX=0x02   MIN=0x03

So a reducing collective maps op (an ALU_OP) onto a CCE/DGE compute-op for the in-transfer reduction. The CCE in-transfer reduce is the AllReduce reduce leg — see CCE In-Transfer Reduce, where kbin_cce_op_to_sdma_cce_op @0x2664c0 performs the op→SDMA-CCE remap and the final all-reduce leg lands as a single CCE packet on Cayman and later. (MED / INFERRED — the mapping is strongly implied by the matching ADD/MAX/MIN encodings and the kbin_cce_op remap, but no static assert states it for 0xC8 specifically.)

GOTCHA — op only matters for reducing ctypes. For ALL_GATHER, ALL_TO_ALL, PERMUTE, PERMUTE_IMPLICIT, and ALL_TO_ALL_V, the lowered DMA runs with DGE_COMPUTE_OP_NONE (pure copy) and op is ignored. It is consumed only for ALL_REDUCE, REDUCE_SCATTER, PERMUTE_REDUCE, and PERMUTE_REDUCE_IMPLICIT. (MED / INFERRED.)

5.4 · The on-chip leg — SB2SB_COLLECTIVE (0xBF, non-pseudo)

The intra-node leg the TOP_SP sequence ultimately programs is the real (non-pseudo) HW collective SB2SB_COLLECTIVE (0xBF, NEURON_ISA_TPB_S3D3_COLLECTIVE_STRUCT). Note 0xBF = 0b101_11111 — upper three bits 0b101, not 0b110, so it is a genuine hardware instruction, not a pseudo-op. It is parameterized by lnc_size_fmt (LNC1/2/4/8 = which NeuronCores form a local group; comment "LncSize setting for sb2sb collective using q7 iDMA", common.h:814) and num_active_channels. The device kernel that decodes and executes it lives in the Cairo firmware — see SB2SB Remote-Copy Collective Kernel and S3D3 Collective (SB2SB, 0xBF). (HIGH / OBSERVED — the microcode binary libnrtucode_internal.so carries decode_sb2sb_collective, S: SB2SB_Collective, and S: POLL_SEM strings, confirming both legs are decoded on the device while 0xC8 is not.)

6 · Position in the collective pseudo-op family

All struct↔opcode bindings from instruction_mapping.json; all sizes probe-verified 64 B unless noted. Every entry below has opcode upper-3-bits 0b110 (pseudo) except 0xBF (0b101, real HW).

NOTE — v2 exposes what v1 leaves to defaults. The two-instruction v2 form (0xD9 + 0xDA, always back-to-back per its header) surfaces channel_id, stream_id, the CC scratch buffer (cc_buf_*), and DMA priority_class — the knobs the v1 single-word 0xC8 lets the runtime pick. The dma_configs.priority_class is a 3-bit field (NEURON_ISA_TPB_DMA_CONFIGS.priority_class : 3, common.h:714); the host maps it to the 5-value priority class P0..P4. Full detail in TriggerCollective2 + Ext. (HIGH / OBSERVED — v2/ext field offsets confirmed in ..._collective2_ext.h:58-90.)

Supporting glue ops a lowered collective sequence draws on: 0xDB PSEUDO_CUR_PROCESSING_RANK_ID (rank loop, §4), 0xD8 PSEUDO_CORE_BARRIER, 0xD5 PSEUDO_SYNC_BARRIER, 0xC1 PSEUDO_DMATRIGGER (§5.1), 0xC3 PSEUDO_DMABARRIER. (HIGH / OBSERVED — opcode enum common.h:263-293.)

7 · Reimplementation checklist

To emit and lower a 0xC8 on a Vision-Q7-compatible engine:

1. Encode the 64-byte word exactly per §1 — header.opcode=0xC8, inst_word_len per ISA, then events, op, dtype, group_id, input/output_tensor_id, num_elements, ctype @ off 32, zeroed reserved1[15], src/dst_offset_elems @ off 48/56. The 0xC8 upper-3-bits 0b110 invariant must hold or NRT will not treat it as a pseudo.
2. Resolve input_tensor_id/output_tensor_id (compiler handles) to NeuronAddrs at model load; apply src/dst_offset_elems * sizeof(dtype).
3. Pick topology from group_id → replica group; supply channel_id/ring defaults (v1 has none — choose them in the rewriter).
4. Compose the TOP_SP (engine 5) program per replica-group iteration: PollSem(0xB3) for the min/decrement semaphore sync (§5.2) + DmaTrigger(0xC1) writing the DMA tail-pointer CSR (§5.1), with barrier/rank-id glue.
5. Set the in-DMA reduce for reducing ctypes: map op (ALU_OP) → CCE_OP (ADD=0/MAX=2/MIN=3) or DGE_COMPUTE_OP; use NONE for non-reducing kinds.
6. Kick each DmaTrigger as a Write to the named queue's tail pointer; the DMA engine moves and reduces the bytes (intra-node leg = SB2SB 0xBF).

For the full host trace (NEFF → load-time rewrite → TOP_SP cc_op SPAD → ring/mesh/hier step → CCE leg), follow A Collective, End to End and TOP_SP Collective Lowering.

See also

• TriggerCollective2 + Ext (0xD9/0xDA) — the SBUF-tensor, two-instruction v2 with channel/stream/cc-buffer/priority knobs.
• ALL_REDUCE — the 0xC7 fixed-form subset and the AllReduce data path.
• TOP_SP Collective Lowering — how 0xC8 becomes a TOP_SP cc_op SPAD program.
• Collective-Type + cc_op Enum Reference — full COLLECTIVE_TYPE, ALU_OP, CCE_OP, DGE_COMPUTE_OP tables.
• S3D3 Collective (SB2SB, 0xBF) — the real-HW intra-node leg.
• PseudoCurProcessingRankID (0xDB) — the replica-group / channel_id / iteration_id vocabulary.
• CCE In-Transfer Reduce — the in-DMA reduce that realizes op.
• SB2SB Remote-Copy Collective Kernel — the device firmware that executes 0xBF.
• A Collective, End to End — orientation trace tying it all together.