The Unified Collective-Communication Architecture

What this page is. The capstone of the collectives chapter — a navigation + mental-model layer that ties the thirteen concrete op pages into one end-to-end stack: compiler pseudo-op → host NRT lowering (SELECT → COMPOSE → EMIT) → TOP_SP on-device sequencer → NCFW ring/mesh/hier algorithm → SB2SB/SDMA data plane → EVT_SEM barrier/semaphore completion. It does not re-derive any sibling page; its value is the spine diagram, the per-boundary handoff table, the cross-links into every Part-10 page, and the open-items ledger. Each consolidated fact cites its source page; the spine HIGH facts additionally carry an addr/offset/symbol re-confirmed against the shipped binaries this session (§0.2).

Provenance. Every structural fact below is established on a committed sibling page (cited inline). The spine addresses/opcodes/enums were re-confirmed directly against the shipped artifacts — the cayman ISA headers (aws-neuronx-gpsimd-customop-lib 0.21.2.0), the host runtime libnrt.so.2.31.24.0 (BuildID 8bb57aba…102e, x86-64, DWARF, not stripped; .text/.rodata/.data all VMA == file-offset, no delta), the NCFW JSON pretty-printer libncfw.so, and the device ucode libnrtucode_extisa.so. All output reads as derived from binary/static analysis only. Confidence is tagged HIGH/MED/LOW × OBSERVED (read from bytes/disasm/header/DWARF this session) / INFERRED (architectural reading over OBSERVED facts) / CARRIED (consolidated from a sibling page at its stated confidence — never inflated).

0 · The architecture in one paragraph + the spine spot-checks

0.1 · The stack in one paragraph

The GPSIMD/NeuronCore collective stack is a six-layer pipeline, with a named artifact handed across each boundary. (1) The compiler emits a 64-byte pseudo instruction targeted at the TOP_SP engine (engine 5) — opcode upper three bits 0b110: TRIGGER_COLLECTIVE 0xC8, TRIGGER_ALL_REDUCE 0xC7, TRIGGER_COLLECTIVE2+_ext 0xD9/0xDA, SENDRECV 0xCB, the barrier trio 0xD8/0xD5/0xC3, CUR_PROCESSING_RANK_ID 0xDB — which never executes on hardware. (2) The host NRT runtime (libnrt.so) decodes + lowers it at model load: it re-encodes the ISA ctype to an enc_op_type, the reduce ALU_OP to an SDMA_CCETYPE, the dtype verbatim to SDMA_DTYPE, SELECTs an enc_alg_type per leg (__select_algorithms @0x14c620 probing enc_can_post_*_operation), COMPOSEs the SDMA step primitives, and EMITs a SPAD cc_op program (create_spad_ctrl_entry @0x232cd0) + DMA descriptor packets (add_dma_packet_cce @0x2307d0) that it DMA-loads onto the device. (3) The TOP_SP — an embedded Xtensa-NX core (engine 5, one per NeuronCore) — is the on-device sequencer: kicked once by a value-1 write to the host_trigger CSR (LOCAL_REG + 0x15a0 = 0x615a0 cayman/mariana, 0x60848 sunda) through the kaena_khal vtable slot +0x708, it walks the cc_op program, per step incrementing EVT_SEM semaphores (the SEMAPHORE_INC window at sp_base + 0x1800, the same window the DMA legs bump) and posting completion as the elected leader. (4) The NCFW management firmware's per-TOP_SP context (rooted at the ncfw_ctx_top_sp key) is that program; the 4-bit algo_type nibble of each cc_op entry routes it to the resident ring, mesh, or hierarchical algorithm. (5) The data plane is the SB2SB collective kernel (S3D3_COLLECTIVE 0xBF, real HW on the POOL/Q7 iDMA): per leg it decodes → cross-die pre-syncs → maps the remote window → rdma_desc_gen (op 8) → rdma_desc_start (op 9, the TX/RX tail-pointer doorbell), moving SBUF → remote-SBUF and, on reducing legs, folding via the SDMA CCE descriptor. (6) Completion/sync is the EVT_SEM 256-semaphore array (arrive INC +0x1800 / wait READ-GE +0x1000 / reset DEC +0x1C00): the per-step peer-semaphore handshake sequences steps within a phase, the NCFW 4-step counted barrier brackets each phase, and the TOP_SP timestamp_inc tick keeps the dies' clocks aligned (tsync). Cross-die addressing rides the 58-bit SoC bitfield over the io_d2d fabric. The transport is bespoke (AL/UDMA hw_exec_queue + EVT_SEM/NQ) — there is no Cadence XRP (§XRP). The one standing boundary: the host→device-core handoff crosses the un-disassemblable Xtensa-LX NCFW core, so the on-core schedule itself is the MED/LOW frontier (§9). [synthesis HIGH/CITED across all anchors; the on-NCFW-core schedule is MED — §9, O-1.]

0.2 · The spine spot-checks (re-confirmed against the binaries this session)

These are the facts the whole synthesis hangs on — re-read directly, not paraphrased:

⇒ Every opcode / address / enum the synthesis depends on is re-grounded; the layers below the binary surface (the on-NCFW-core schedule) are the consolidated MED/LOW findings of the source pages, carried forward at their stated confidence.

1 · The layered stack — the handoff at each boundary

The collective is a vertical pipeline; at each boundary the producer hands the consumer a named artifact. This is the integrating spine — read top-to-bottom.

1.1 · The spine diagram

  ┌────────────────────────────────────────────────────────────────────────────┐
  │ L1  COMPILER (Neuron/XLA backend)                                            │
  │   emits a 64-B PSEUDO word, opcode upper-3-bits 0b110, targeted at TOP_SP #5 │
  │   0xC7 / 0xC8 / 0xD9+0xDA / 0xCB / 0xD8 / 0xD5 / 0xC3 / 0xDB   (never on HW)  │
  └──────── NEFF instr stream + def.json replica_groups/dma_queues ─────────────┘
                                   │
  ┌────────────────────────────────▼───────────────────────────────────────────┐
  │ L2  HOST NRT  (libnrt.so, nrt_collectives.cpp)   — at MODEL LOAD             │
  │   DECODE ctype/op/dtype  →  re-encode enc_op_type / SDMA_CCETYPE / SDMA_DTYPE│
  │   SELECT  __select_algorithms @0x14c620  → enc_can_post_*  → enc_alg_type    │
  │   COMPOSE recv_reduce_* / direct_recv_*  (per-leg SDMA step primitives)      │
  │   EMIT    create_spad_ctrl_entry @0x232cd0 (cc_op table) +                   │
  │           add_dma_packet_cce @0x2307d0 (DMA desc packets)                    │
  └──── (a) SPAD cc_op program (HBM → DMA) ── (b) host_trigger doorbell 0x615a0 ─┘
                                   │  vtable +0x708 (set_host_trigger)
  ┌────────────────────────────────▼───────────────────────────────────────────┐
  │ L3  TOP_SP on-device SEQUENCER  (Xtensa-NX, engine 5, per NeuronCore)        │
  │   kicked once → WALKS the cc_op program; per step issues DMA tail-ptr INCs   │
  │   into EVT_SEM SEMAPHORE_INC (sp_base + 0x1800 + idx*4) and EVT_SEM wait/dec │
  └──────── the NCFW per-TOP_SP context IS this program (algo_type nibble) ──────┘
                                   │
  ┌────────────────────────────────▼───────────────────────────────────────────┐
  │ L4  NCFW ALGORITHM  (libncfw ctx + Xtensa-LX firmware images)               │
  │   cc_op byte0[0:3]=algo_type → dispatch (DRAM+0xB0 12-entry jump table)      │
  │   route → RING (NCFW-04) | MESH (NCFW-05) | HIER (NCFW-06);  emit SB2SB legs │
  └──────── per channel: next/prev neigh + recv/send/post/dma_compl sema addrs ──┘
                                   │
  ┌────────────────────────────────▼───────────────────────────────────────────┐
  │ L5  DATA PLANE  (SB2SB collective kernel, 0xBF S3D3_COLLECTIVE, POOL/Q7 iDMA)│
  │   decode → cross-die PRE-SYNC → program_window → rdma_desc_gen (op8) →       │
  │   rdma_desc_start (op9): TX writes M2S TDRTP_inc / RX writes S2M RDRTP_inc;  │
  │   reduce legs FOLD via the SDMA CCE descriptor                              │
  └──────── local_sem (src release) + peer recv_sema (dst ready) ───────────────┘
                                   │
  ┌────────────────────────────────▼───────────────────────────────────────────┐
  │ L6  COMPLETION / BARRIER / SEMAPHORE  (EVT_SEM 256-array; NCFW barrier; tsync)│
  │   step↔step  = per-peer kring/ring/mesh handshake  (INC +0x1800 / GE +0x1000)│
  │   phase bracket = NCFW 4-step counted barrier (≤4 leaders, steps {1,4,4,1})  │
  │   clock align  = TOP_SP timestamp_inc tick → dev tsync                       │
  │   to host      = reserved-sema INC and/or NQ entry (stride 0xa0)             │
  └──────────────────────────────────────────────────────────────────────────────┘

1.2 · The per-boundary handoff table

NOTE — the two-level trigger model. L2's host_trigger doorbell (top-sp-lowering §3a) starts the TOP_SP program once per executable; L3's per-step DMA-tail INC into EVT_SEM +0x1800 (top-sp-lowering §3b) is the per-iteration hand-off the DMA legs use. The program, once started, spins on the same +0x1800 window the DMA legs increment. Two distinct surfaces, one semaphore array.

2 · The pseudo-op catalog — the full opcode roster

All opcodes spot-verified in cayman common.h this session (§0.2 #1). Each is a 64-byte TPB instruction word; the pseudo class = opcode upper-three-bits 0b110 (0xC0..0xDF). 0xBF (below 0xC0) and 0xF1 (EXTISA) are real HW, not pseudo. "Triggers algo" = which L4 path the lowering drives.

GOTCHA — the DMABARRIER token mismatch. The opcode enum constant is the no-underscore OPCODE_PSEUDO_DMABARRIER (=0xc3) but the struct/JSON-key are the underscore PSEUDO_DMA_BARRIER; instruction_mapping.json's opcode TOKEN is a dangling stale Tonga-era underscore form. Same single 0xC3 (value 195); value-based decoders are unaffected, token-resolving tools see a dangling binding. CORE/SYNC have no such mismatch. [CARRIED — dma-barrier §6, HIGH/OBSERVED] See dma-barrier, collective-enums.

NOTE — ALL_REDUCE has three carriers, not one. It is both a distinct opcode (0xC7) and a ctype value (0x1) on 0xC8 and on 0xD9/0xDA. The reimplementer must accept all three forms. [CARRIED — all-reduce §1, HIGH/OBSERVED]

3 · The algorithm family — enc_alg_type, selection, reduce carriers

3.1 · The enc_alg_type selector (11 values, spot-verified §0.2 #5)

enc_alg_type is exactly the low nibble of the firmware cc_op algo_type field. The 11 values, the family they belong to, and the L4 algorithm page each routes to:

enc_alg_mesh_type (= cc_op algo_sub_type, 3-bit): FULL_MESH0 / GROUPED_MESH1 / MESH_TRN2 2 / MESH_SWITCH3 / MESH_INVALID4. [values HIGH/OBSERVED §0.2; the host↔firmware identity is MED/INFERRED — §8 C-7.] Full enum reconciliation: collective-enums.

3.2 · When each is selected — __select_algorithms @0x14c620

__select_algorithms (spot-verified §0.2 #6) assigns one enc_alg_type per leg by probing the enc_can_post_*_operation predicates. A hierarchical all-reduce carries up to five per-leg choices. The legality matrix (from the libnrt .rodata assertion strings):

Gates: SINGLE_CYCLE_RING is ALL-REDUCE-only (op_type == ENC_ALLREDUCE guard, in enc_can_post_single_cycle_ring @0xfaa80); KANGARING is intra-only (enc_can_post_kangaring @0xfbe70 rejects inter comms). The exact numeric algo a given (op, scope, world-size, dtype, topology) resolves to depends on the enc_can_post_* thresholds, which are not enumerated. [HIGH/OBSERVED gates; threshold→numeric is MED — §9 O-3.]

3.3 · The algorithm decompositions

• RING all-reduce = reduce-scatter (N-1 recv_reduce_* steps, CCE-fold) + all-gather (N-1 direct_recv/send copy steps), over directed channel rings; flow control = ring_ctx recv_cnt/send_credit/run_state. ring-kangaring, all-reduce.
• KANGARING = the same ring tape but reduction_type_t = KANGARING_NR1W(2): the primary rank reads-and-reduces from a vector of N peers (kring_peer_semas) and writes to 1 next; latency-opt for intra. ring-kangaring.
• MESH all-reduce = a flat event tape (50 events SUNDA / 108 the rest); each event = wait event_wait_sema ≥ wait_val → fire ≤2 dma_apb_bcast → INC ≤3 direct_trigger_sema at neighbour dies; the SoC CAYMAN_ID/DIE bits route the d2d hop. mesh-collective.
• HIERARCHICAL all-reduce = intra reduce-scatter (RING) → inter all-reduce (MESH/RING/RDH) → intra all-gather (RING), sequenced by a single run_state byte; the hier config is an 8-byte handle to a firmware-resident level table — it reuses the co-resident ring+mesh descriptors, does not embed them. hierarchical-collective.
• SENDRECV = the irreducible building block: a single directed SB2SB leg, no chaining, no fold. enc_metaring_primitive::__compose_channel_ring calls __compose_p2p_channel — a ring step is literally a sendrecv. sendrecv.

GOTCHA — pring is not a value. There is no PRING enc_alg_type (the 11-enum has none). The enc_metaring "metaring" family + __compose_pipeline_allreduce are the nearest landed artifacts; a dedicated pring decode is OPEN (§9 O-2). The stub pring-descriptors holds what is known.

3.4 · The reduce-op + dtype carriers

• reduce op rides ISA op(ALU_OP)@12 → SDMA_CCETYPE {ADD0,FMA1,MAX2,MIN3,EXT4,GCE5} (== ISA CCE_OP {ADD0,MAX2,MIN3} byte-superset; spot-verified §0.2 #4/#7). It is realized in the SDMA CCE descriptor (add_dma_packet_cce @0x2307d0) — not in the 0xBF word and not in the cc_op spad word. all-reduce, collective-enums, cce-in-transfer.
• reduction_type_t (the fold pattern): RING_2R1W=0, RING_2R2W=1, KANGARING_NR1W=2. ring-kangaring.
• dtype rides dtype@13 → SDMA_DTYPE (byte-identical pass-through). The CCE-reducible set = {BF16, FP16, FP32R, FP8_E3/E4/E5} (plain FP32 and integers absent). all-reduce, collective-enums.

4 · The barrier / semaphore model

4.1 · The EVT_SEM substrate (spot-verified §0.2 #8)

A 1 MiB SoC aperture per instance: 256 events (1-bit) + 256 semaphores (32-bit), with op-aliased windows over the same 256 physical counters:

The TOP_SP-hosted copy (cayman_get_sp_sema_{r,s,i,d}_ofst) is the same four windows on the per-TOP_SP base 0x8280000000 (stride 0x40000000, <<0x2f die bit). [HIGH/OBSERVED; per-page detail: top-sp-lowering §4.] top-sp-lowering, core-barrier.

4.2 · The three-barrier trilogy (the scope ladder)

NOTE — no timeout in any of the three. The participant count N (CORE_BARRIER semaphore_value) = the pcores of that VNC; all three spin unconditionally — liveness is a compiler/runtime guarantee, not a hardware timeout. [CARRIED — core/sync/dma-barrier §8, HIGH/OBSERVED]

4.3 · EVT_SEM + the NCFW 4-step barrier (the COARSE cross-rank fence)

The NCFW device barrier is a 4-step counted-semaphore barrier that brackets each ring/mesh phase: barrier_step_config[4] (52 B each), barrier_step_sizes {1,4,4,1}, op {0:SIGNAL/arrive, 1:WAIT-GE-and-dec, 2:NOP}. Canonical:

  step0 local ARRIVE → step1 SIGNAL ≤4 leader peers → step2 WAIT on those 4 → step3 local RELEASE
                                          ("4" = NEFF_DEVICE_BARRIER_MAX_LEADERS)

The host barrier (16 B barrier_start@0 / barrier_done@8) merges with the device done sema in host-CC mode; host-side timeouts exist there ("non-leader workers waiting at barrier exceeded %d seconds"). Named TOP_SP sema-ids: ENCD_TRIGGER=0, COMPLETION=1, BARRIER_SEMA0/1=2/3, BARRIER_DMA_SYNC=4, FUNCTION_SWITCH=5, HOST_BARRIER0/1=6/7. neff-device-barrier, neff-host-barrier.

4.4 · Per-step handshake vs global barrier — they COMPOSE

The per-step kring/ring/mesh peer handshake (mine + peers[0..2] / recv_sema / send_sema) is a per-chunk inc+wait-GE that sequences steps within a phase. The NCFW 4-step barrier is the global fence that brackets the phase. Both hit the same EVT_SEM ops (add_semaphore_inc 0x10A0/21 = INC +0x1800; add_semaphore_wait_ge_and_dec 0x10A0/20 = READ-GE +0x1000 then DEC +0x1C00). [CARRIED — ring-kangaring §3/§7, neff-device-barrier; HIGH ops / MED schedule]

4.5 · tsync (the clock substrate under the barriers)

The TOP_SP timestamp_inc tick (top_sp_ram +0x4, RW[23:0] reset 0x400) advances a global timestamp; the dev tsync struct (timestamp_local / timestamp_tpb / timestamp_tpb_val) is the firmware's read/align, fanned via the EVT_SEM/APB-broadcast fabric. Barrier sequences ops; tsync keeps the dies' clocks aligned. [doorbell HIGH; the timestamp→tsync align is MED — runs on the LX/NX cores, §9.] spad-ccop-tsync, top-sp-lowering.

5 · The data plane — SB2SB 0xBF, RDMA desc gen/start, M2S/S2M roles

5.1 · SB2SB / S3D3_COLLECTIVE (0xBF, real HW, POOL)

The 64-B word: events@4, in_dtype@12/out_dtype@13 (validator enforces in==out, a constrained-identity cast slot), src_mem_pattern TENSOR3D@16 (ADDR4 base + 3 int16 strides + 3 uint16 dims, ≤256 elements, SBUF-only), lnc_size_fmt@32 (LNC1 self / LNC2 pair; LNC4/LNC8 rejected → intra-node ≤2-NC), num_active_channels@33 (1..128), dst_mem_pattern TENSOR3D@48. The 0xBF word has no reduce-op field — the reduce lives in the SDMA CCE descriptor on reduce-copy legs. S3D3 does both structured COPY (all-gather legs, add_dma_packet) and REDUCE-COPY (RS/AR legs, add_dma_packet_cce). Two co-operating engines: SEQ control (0xBF handler) + POOL/Q7 data plane (moves the bytes). s3d3-collective, cce-in-transfer.

5.2 · rdma_desc_gen (op 8) → rdma_desc_start (op 9)

• rdma_desc_gen (op 8) reads {local_sem, remote_sem, remote_core_id, remote_routing_id, dma_engine_mask, src/dst addr, free_dim_bytes} and builds, per selected DMA engine, an SDMA CME BD ring (16-byte BD: length+gen-tag, CME COPY cmd, buf_ptr u64 SoC addr) copying [src, src+free_dim_bytes] across 128 SBUF partitions, plus a LOCAL sema descriptor (bumps local_sem = source release) and a REMOTE sema descriptor (bumps remote_sem on the peer = dst ready). Can be issued early (off the critical path).
• rdma_desc_start (op 9) does DRAIN → role split (PRID parity) → TX/RX doorbell.

The M2S/S2M role mapping (pinned from the DRAM "TX\0RX\0"/"M2S\0S2M\0" token block):

Both are s32i.n a3,a4,0 tail-pointer-increment stores = the DmaTrigger primitive. s3d3-collective, sendrecv.

5.3 · Cross-die addressing

The descriptor buf_ptr carries the SoC bitfield: LOCAL[46:0], DIE[47], CAYMAN_ID[53:48] (64-die mesh), CAYMAN_ID_VALID[54] (=1 ⇒ route by chip). remote_routing_id is rewritten into {CAYMAN_ID, CAYMAN_ID_VALID} (+ neighbour EXIT_DIE[51]/NEIGHBOR_ROUTE[52]) to turn a local dst_addr into a remote-die address; the copy traverses the io_d2d fabric. The remote SBUF must first be mapped into the local Q7 32-bit window by program_window. [bitfield HIGH/OBSERVED; routing_id→high-bits rewrite MED.]

5.4 · Completion — semaphore-over-DMA, NO interrupt

Two-semaphore (local_sem source-release / remote_sem data-ready) + a generation-tag busy-poll. No interrupt path for the data plane — completion is semaphore-over-DMA + poll; a descriptor count/role mismatch is a hard error. [HIGH/OBSERVED logs; exact Q7 poll site MED — FLIX-desync, §9.]

6 · The three-reduce-layer placement (for all_reduce)

There are three distinct reduce-op enums (plus the public nrt_op_type API enum and the DGE enum), at distinct fold scopes that compose innermost → outermost. They agree only on the abstract names (ADD/MAX/MIN) — they are separate packed types with different encodings (the single most common reimplementation error). All spot-verified in the cayman header this session (§0.2 #4).

Placement for all_reduce. A hierarchical all-reduce folds at progressively coarser scopes: the intra-die ring reduce-scatter folds via the cross-core SDMA CCE (SDMA_CCETYPE), the inter-die mesh leg folds again via the same CCE. Layers (1) and (1.5) are on-chip compute datapaths a kernel may use to produce the per-rank tensor — distinct from the collective fold. The collective op carries the SDMA CCE op only; the cc_op spad word and the 0xBF word do not carry the reduce op.

GOTCHA — alignment islands (do not interchange). REDUCE_OP and CCE_OP/SDMA agree on ADD=0, MAX=2 but disagree on slot 1 (AVERAGE vs FMA) and slot 3 (OR vs MIN). DGE uses a NONE-based numbering. nrt_op_type {ADD0,FMA1,MAX2,MIN3, INVALID15} == SDMA_CCETYPE for {ADD,FMA,MAX,MIN}; translate_op_type is identity. The cayman CCE_OP_ADD line even comments "Same encoding as SDMA CCE op encoding" and annotates Multiply = 0x01 (the SDMA slot-1 the ISA CCE_OP skips). [CARRIED — collective-enums; HIGH/OBSERVED]

7 · Cross-generation architecture notes (full matrix → GEN-08)

The exhaustive per-gen support matrix lives in master-capability-matrix — cite it. Codename binding (authoritative, not inverted): arch_id 0x05 = v2 = SUNDA, 0x0c = v3 = CAYMAN (the reference SoC), 0x14 = v4 = MARIANA, 0x1c = v4+ = MARIANA_PLUS; MAVERICK v5 (coretype 37) is the known-but-NOT-shipped 5th gen. lx-isa-naming-archid-synthesis.

Schema-stable across the 4 shipped gens (SUNDA..MARIANA_PLUS):

• The collective opcodes + their 64-B struct layouts are byte-identical (only the ISA header for NC-vN comment differs). [CARRIED HIGH/OBSERVED]
• The host enc_*/SDMA_* enums are single-binary (one libnrt serves all gens).
• The NCFW cc_op command-word schema + tsync struct + barrier step descriptor (52 B)
  • ring channel + mesh event (80 B) + hier (8 B) layouts are schema-wide byte-identical.
• The EVT_SEM 256-array geometry + four op windows are gen-stable.

Scales / diverges per gen:

NOTE. MARIANA == MARIANA_PLUS at the device-image level (byte-identical EXTISA + NCFW DRAM; only an NCFW-IRAM code-shift differs). MARIANA_PLUS is a feature-flag refresh (NEURON_RT_DBG_V4_PLUS), adds 0 new ISA/EXTISA bytes. The CCE-reducible collective dtype set stays {BF16,FP16,FP32R,FP8_E3/E4/E5} across all gens. [CARRIED HIGH/OBSERVED; v5 interiors INFERRED — §9 O-7.]

QUIRK — SYNC_BARRIER on v5. The "event-based SYNC_BARRIER" reading holds for v2/v3/v4 but not maverick: v5 drops the EVENT family, so a SYNC_BARRIER there must lower via the SEM/NONBLOCKING forms. The struct is byte-identical; the lowering target differs. [CARRIED — sync-barrier, collective-enums §9.1, HIGH/OBSERVED]

8 · Residual contradictions reconciled

These were divergences among the source findings; resolved (binary wins):

CORRECTION carried for the Part-10 reconcile (C-4, naming both pages). all-reduce's summary table row B describes the host-trigger as going "via kaena_khal vtable +0x710". That shorthand is imprecise: the trigger write (set_host_trigger) is slot +0x708; +0x710 is only the offset getter (re-verified §0.2 #3). All-reduce's prose is already correct (it calls +0x710 "the offset getter; the adjacent set_host_trigger performs the write"); only the table cell uses the loose form. top-sp-lowering is the authoritative split and already flags this. No re-derivation contradiction — a documented, already-CORRECTED divergence.

9 · The open-items ledger

10 · Confidence summary

• HIGH / OBSERVED or SPOT (re-verified §0.2): all collective opcodes + COLLECTIVE_TYPE/CCE_OP/DGE_COMPUTE_OP/LNC_SIZE_FMT enums; enc_alg_type 0..10 + SDMA_CCETYPE; host_trigger 0x615a0/0x60848 + the vtable +0x708/+0x710 split; EVT_SEM windows +0x1000/1400/1800/1C00; the NRT lowering symbol addresses (__select_algorithms @0x14c620, add_dma_packet_cce @0x2307d0, create_spad_ctrl_entry @0x232cd0, the enc_can_post_* family); the device SB2SB strings + ncfw_ctx_top_sp/run_state. The pseudo-op catalog (§2), the barrier model (§4), the data plane + M2S/S2M roles (§5), the algorithm family + legality (§3), the three-reduce reconciliation (§6), the schema-stable + per-gen deltas (§7) — all CITED HIGH at their source-page confidence.
• MED / INFERRED-STRONG: the host enc_alg_type ≡ firmware algo_type identity (C-7); the per-leg numeric algo per workload (O-3); the leg↔NCFW-step correspondence + the per-step schedule; the routing_id→SoC high-bit rewrite; the tsync align algorithm; "TOP_SP executes the NCFW firmware" (O-8).
• LOW / OPEN: the NCFW LX-core case-body decode (O-1); pring (O-2); the on-wire CCE descriptor bitfield (O-4); the runtime-bound spad/kring/barrier slot details (O-5); the control-plane boundaries (O-6); the coretype→silicon binding + MAVERICK arch_id (O-7).

Cross-references — the whole collectives chapter

Pseudo-op catalog (Part-10 ops): TriggerCollective 0xC8 · TriggerCollective2 + Ext 0xD9/0xDA · ALL_REDUCE 0xC7 · S3D3 Collective 0xBF · PseudoCurProcessingRankID 0xDB/0xDC · CORE_BARRIER 0xD8 · SYNC_BARRIER 0xD5 · DMABARRIER 0xC3 · SENDRECV 0xCB · RDMA/DMA Pseudo-Ops 0xb8/0xbb/0xf1 · Collective-Type + cc_op Enum Reference · TOP_SP Collective Lowering · XRP Host↔DSP Messaging (the bespoke transport verdict)

NCFW algorithm layer: NCFW Main Dispatch Loop · Ring + Kangaring · Mesh Collective · Hierarchical Collective · NCFW IRAM Images · LX ISA Naming / arch_id Synthesis

Data plane / reduce / cross-gen / orientation: CCE In-Transfer Compute (the cross-core reduce) · Master Capability Matrix (cross-gen support) · A Collective, End to End (the orientation trace)