PSEUDO_SYNC_BARRIER (0xD5)

NEURON_ISA_TPB_PSEUDO_SYNC_BARRIER_STRUCT, opcode 0xD5 (NEURON_ISA_TPB_OPCODE_PSEUDO_SYNC_BARRIER). This is the within-core, all-engine rendezvous of the Cairo TPB ISA: a single 64-byte pseudo instruction the compiler emits at a loop boundary to align the execution engines of one NeuronCore (PE / ACT / POOL / DVE / TPB_SP, plus TOP_SP as orchestrator) before the next iteration. It is the second member of the barrier trilogy and the finest-grained of the three:

The header doc comment is the anchor, verbatim from aws_neuron_isa_tpb_pseudo_sync_barrier.h:18-20:

"A struct to be used with PSEUDO_SYNC_BARRIER opcode. Used as all-engine-barrier for loop-on-chip flow."

Every struct/enum/opcode/offset fact on this page is 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), the cayman register-schema JSON (tpb_events_semaphores_axi.json), the instruction_mapping.json struct→opcode table, and the device microcode libnrtucode_internal.so. All are binary-derived artifacts of the package and are cited directly. Struct layouts were re-checked with a self-contained gcc probe printing sizeof/offsetof for every field.

NOTE — pseudo-op, never executed as-is. 0xD5 lives in the pseudo range 0xC0..0xDF. The OPCODE enum states the rule verbatim at aws_neuron_isa_tpb_common.h:263 (on the 0xC1 line): "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." 0xD5 = 0b110_10101 ✓. The device-side GPSIMD microcode (libnrtucode_internal.so) has no decoder string for any barrier — it decodes the lowered legs (Event_Semaphore, POLL_SEM), confirming the rewrite happens host-side. (HIGH / OBSERVED.)

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

aws_neuron_isa_tpb_pseudo_sync_barrier.h:25-29. Natural C alignment applies; total size is 64 bytes = one TPB instruction word.

typedef struct NEURON_ISA_TPB_PSEUDO_SYNC_BARRIER_STRUCT {
    NEURON_ISA_TPB_HEADER header;        // 4   ( 0 -  3)  opcode=0xD5
    NEURON_ISA_TPB_EVENTS events;        // 8   ( 4 - 11)  inline wait+update block
    uint8_t               reserved0[52]; // 52  (12 - 63)  padding — NO id/sema/metadata
} NEURON_ISA_TPB_PSEUDO_SYNC_BARRIER_STRUCT;

gcc probe over the real shipped headers (-I…/neuron_cayman_arch_isa/tpb, offsetof/sizeof):

Probe output (independently reproduced): sizeof == 64, off events == 4, off reserved0 == 12. The NEURON_ISA_TPB_OPCODE enum is NEURON_ISA_PACKED (1 byte; probe sizeof(NEURON_ISA_TPB_OPCODE) == 1), so header.opcode occupies exactly byte 0. (HIGH / OBSERVED.)

The decisive structural fact: SYNC_BARRIER carries the common header + the inline events block and nothing else. There is no barrier identifier, no semaphore handle, no count, and no mask — only reserved0[52] to the 64-byte word. This events-only shape is the encoding fingerprint of the binary-flag (event) sync mechanism (§3, §6).

NOTE — no structural validator. aws_neuron_isa_tpb_assert.h:13565 contains only the bare section marker //* pseudo_sync_barrier_assert.h between empty //***** blocks — no validator body. The neighbouring pseudo_core_barrier_assert.h (:12648) and pseudo_dma_barrier_assert.h (:12654) markers are likewise empty. Contrast ctrl_es.h:55 (the lowering target), which carries a full ISA_STATIC_ASSERT and an is_valid_ctrl_es predicate. Barriers are runtime-validated, not statically asserted. (HIGH / OBSERVED.)

1.1 · Binding (mapping + union)

• instruction_mapping.json:175-176: "NEURON_ISA_TPB_PSEUDO_SYNC_BARRIER_STRUCT" → ["NEURON_ISA_TPB_OPCODE_PSEUDO_SYNC_BARRIER"]. The mapping token spells the opcode name exactly as the enum constant.
• aws_neuron_isa_tpb_util.h:66: INST_UNION member NEURON_ISA_TPB_PSEUDO_SYNC_BARRIER_STRUCT pseudo_sync_barrier;.
• aws_neuron_isa_tpb_common.h:283: NEURON_ISA_TPB_OPCODE_PSEUDO_SYNC_BARRIER = 0xd5.
• aws_neuron_isa_tpb_enums.h:325: member of is_valid_enum_opcode (a valid TPB opcode). It is not a member of is_valid_enum_pseudo_opcode (enums.h:352-356), which is the unrelated tiny set {INVALID, PSEUDO_EXIT_EXECUTION, PSEUDO_LIBRARY_RELOAD_INDEX}. (HIGH / OBSERVED.)

CORRECTION — SYNC is clean; the naming anomaly is DMA's alone. The DMA sibling's enum constant is NEURON_ISA_TPB_OPCODE_PSEUDO_DMABARRIER (no underscore, common.h:265), but instruction_mapping.json:109 references NEURON_ISA_TPB_OPCODE_PSEUDO_DMA_BARRIER (with underscore) — a token that does not exist in common.h (rg -c → 0 hits in the header, 1 hit in the mapping). SYNC_BARRIER has no such mismatch: OPCODE_PSEUDO_SYNC_BARRIER resolves in both files (rg -c → 1, 1). When wiring a decoder, key 0xD5/0xD8 off the enum value, not the mapping string; for 0xC3 the mapping string is a dangling reference. (HIGH / OBSERVED.)

1.2 · Cross-arch invariance

The struct body is byte-identical on cayman / mariana / maverick / sunda — diff of the four aws_neuron_isa_tpb_pseudo_sync_barrier.h files (skipping the banner) is empty; only the ISA header for NC-vN comment line differs (mariana = NC-v4, maverick = NC-v5, sunda = NC-v2, cayman = NC-v3). Opcode 0xD5, the mapping binding, and the union member are identical on all four. (HIGH / OBSERVED for the four ISA-header families.)

NOTE — v5/MAVERICK interiors are inferred. The ISA struct exists for maverick, but maverick ships a different register-schema format: there is no tpb_events_semaphores_axi.json under arch-headers/maverick/. Instead the event/semaphore CSR appears as tpb_coll_sync.h with offsets such as AWS_REG_MAVERICK_TPB_COLL_SYNC_SEMAPHORE_SUBTRACT_OFFSET = 0x10000 and …WRITE_IF_ZERO_OFFSET = 0x18000. The cayman/mariana/sunda window map (§3.1) is byte-verified; the v5 event-window placement is header-OBSERVED only — treat any v5 offset as INFERRED. (MED / INFERRED.)

2 · The events block (the sync primitive)

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

The header even ships the Default impl in a comment (common.h:427-429): Events{wait_mode: None, wait_idx: 0, update_mode: None, update_idx: 0, semaphore_value: 0}.

The 8-bit wait_idx/update_idx index the 256-entry EVT/SEM array exactly (2^8 == 256, §3.1). This same inline events block is carried by every TPB instruction that can synchronize — the barriers, EVENT_SEMAPHORE (ctrl_es.h:48), POLL_SEM (ctrl_poll_sem.h:45), TriggerCollective. What distinguishes SYNC_BARRIER from CORE_BARRIER is which mode family the lowered block selects (event vs semaphore), not the field types. (HIGH / OBSERVED.)

3 · The events-based mechanism (events vs semaphores)

This is the central decode. The WAIT_MODE and UPDATE_MODE enums each carry two disjoint families — an EVENT family (1-bit flag) and a SEMAPHORE family (32-bit counter) — and the SYNC-vs-CORE distinction is which family the lowered barrier uses.

NEURON_ISA_TPB_WAIT_MODE (common.h:313-330):

NEURON_ISA_TPB_UPDATE_MODE / "Event Update Mode" (common.h:333-369; the arrive side). _READ fires after reads complete, _COMPLETE (+0x10) after the instruction is globally observed (common.h:333-343):

So SYNC_BARRIER lowers to event modes (WAIT_FOR_EVT_SET = 0x6 / EVT_SET_COMPLETE = 0x11); CORE_BARRIER lowers to semaphore modes (WAIT_FOR_SEM_GE_IMM = 0x5 / SEM_INC_COMPLETE = 0x13). (WAIT/UPDATE enum values HIGH / OBSERVED; the SYNC→event, CORE→sem family assignment is the architectural reading — MED / INFERRED, grounded in the events-only struct + the two HW window families below.)

3.1 · The HW EVT/SEM window map (cayman register schema)

One physical block holds both primitives. From cayman-arch-regs/csrs/tpb/tpb_events_semaphores_axi.json (RegFile.RegistersBundleArrays, AddrWidth 20 / DataWidth 32), re-verified byte-identical in the sunda (v2) copy:

(BundleSizeInBytes = 0x4 per entry, ArraySize = 256 per window.)

• SYNC_BARRIER (events) maps to the tpb_events window @0x0 — 256 one-bit flags (set = 1 / clear = 0), addressed RW. Arrive = an EVT_SET write of 1 to tpb_events[idx]; wait = a poll of tpb_events[idx] for == 1 (WAIT_FOR_EVT_SET), optionally consuming it with _THEN_CLEAR (0x7). The events.semaphore_value u32 is unused (a flag has no count).
• CORE_BARRIER (semaphores) maps to the tpb_semaphores_* windows @0x1000.. — 256 32-bit counters: arrive = INC (+0x1800), wait = read-GE (+0x1000), reset = DEC (+0x1C00).

The semantic split (why the struct payload differs). A 1-bit event is a binary flag — "engine X reached the barrier point" — set once, polled by peers, optionally cleared. It is the right primitive for a small fixed set of participants (the ≤6 engines of one core), where each engine owns its own event bit; no counting is needed. A 32-bit semaphore counts arrivals against a threshold N (the GE-wait counted barrier) — the right primitive when the participant count is data-dependent or larger (the pcores of a VNC). That is why SYNC has no semaphore/value/id and CORE does: the mechanism dictates the payload. (window/enum facts OBSERVED; the mechanism→payload reading is INFERRED-STRONG / MED.)

NOTE — egress as a notification, not just a busy-poll. An event SET can surface as an async notification-queue record instead of an on-chip spin, gated by events_semaphores.notific_ctrl.notifications_en ("Enable events/semaphores notifications on writes", tpb.json:1106-1110, notific_ctrl at AddressOffset 0x8). The lowered SYNC arrive can therefore be polled on-chip or raise a notification. (window/notif gate OBSERVED; that SYNC's lowering uses it is INFERRED / MED.)

4 · The participant model — all engines of one core

The doc word is "all-engine-barrier". The participant set is the NEURON_ISA_TPB_NEURON_ENGINE enum (common.h:139-146, section header "Execution Engines"):

So SYNC_BARRIER aligns the execution engines of one NeuronCore: the five compute/data engines PE/ACT/POOL/DVE/TPB_SP, plus TOP_SP as the sync-orchestration engine — the 5/6-engine model. (enum HIGH / OBSERVED; the "barrier spans these engines" reading INFERRED-STRONG from the doc word "all-engine".)

How the barrier knows which engines participate. SYNC_BARRIER carries no participant bitmask and no count — the struct is events-only (§1). The participant set is carried implicitly as the per-engine EVENT bits the runtime wires up: each arriving engine is assigned an event bit (or a producer/consumer event pair) in the 256-entry tpb_events array; the runtime emits, per engine, an EVT_SET (arrive) on that engine's bit and a WAIT_FOR_EVT_SET (wait) on the peers' bits. The single inline events block seeds one wait/set pair; the runtime fans it into the per-engine set. The coarse cross-rank barrier (NCFW) does carry an explicit dma_engines_bitmap — but that is a different, broader mechanism (§6); the within-core SYNC barrier needs no such field because one core's engine set is fixed and small. (absence of a mask/count is HIGH / OBSERVED; the per-engine event-bit fan-out is INFERRED / MED — the exact runtime allocation lives in NRT/KRT, not the headers.)

The participants are within one core — not cross-core. The cross-core (pcores of a VNC) scope is CORE_BARRIER's (0xD8, counted semaphore); the cross-rank/die scope is the NCFW collective barrier. The 8-core SPMD context in which these cores execute is described in The Multicore API (8-core SPMD) and The 8-Core SPMD Execution Model + Teardown.

5 · The device lowering (0xD5 → EVENT_SEMAPHORE 0xA0)

There is no dedicated firmware handler for opcode 0xD5 — the expected result for a pseudo-op. The argument is three-legged:

(a) 0xD5 is pseudo (§1). Pseudo-ops are rewritten by NRT/KRT into non-pseudo HW instructions before the stream reaches the engine sequencers, so the device handlers never see 0xD5; they see the lowered form.

(b) SYNC_BARRIER is absent from every binary. A repo-wide search for PSEUDO_SYNC_BARRIER / pseudo_sync_barrier over extracted/ returns hits only in the four ISA-header packages + instruction_mapping.json — zero non-header hits. The device microcode libnrtucode_internal.so's instruction-name printer table (the "S: …" strings) contains "S: Event_Semaphore\n", "S: Event_Semaphore Rng Clr\n", "S: POLL_SEM\n", "S: EngineNop\n", "S: NOP\n", "S: SB2SB_Collective\n" — but no SyncBarrier/CoreBarrier/DmaBarrier string. The interpreter logs "S: Dispatch opcode=0x%x\n" and decodes only the lowered legs. (HIGH / OBSERVED — exhaustive rg + the binary string table.)

(c) The lowering target is EVENT_SEMAPHORE (0xA0). Its op struct is CTRL_ES (ctrl_es.h), bound CTRL_ES_STRUCT → OPCODE_EVENT_SEMAPHORE (instruction_mapping.json:15-16), union member ctrl_es (util.h:22). The doc comment is the close (ctrl_es.h:20-43):

"EventSemaphore … Wait for the 'events' field wait condition (if any) … perform any 'events' wait action (e.g. WaitForEvtSetThenClear) … The instruction can now 'execute' (it's a nop other than the event synchronization) … Set the 'events' field update, if specified … an execution engine would likely launch two 'wait' requests to a centralized event/semaphore block … This instruction does not read any memory sources => all updates must be triggered off '_Complete'."

CTRL_ES is 64 B (ISA_STATIC_ASSERT(sizeof == 64), ctrl_es.h:55; probe confirms): header@0, events@4 (the same 8-B block as the barrier), setter_signature@12, reserved0[19]@13, events_extended@32 (a 12-B NEURON_ISA_TPB_EVENTS_EXTENDED for a second wait/set), reserved1[20]@44. Its validator es_updates_on_completion (ctrl_es.h:72-95) enforces that every update mode be a _Complete form (EvtSet_Complete, EvtClr_Complete, SemInc_Complete, …) — exactly the rule the doc states. A SYNC_BARRIER (events-only) lowers naturally into one — or a per-engine set of — EVENT_SEMAPHORE ops that WAIT_FOR_EVT_SET on peers' event bits and EVT_SET_COMPLETE their own. (the 0xD5 → 0xA0 lowering is MED / INFERRED — EVENT_SEMAPHORE is the only HW op whose entire purpose is event-wait/set "nop other than the event sync", and its events block matches the pseudo-op's byte-for-byte; the exact 1:1 expansion list is not in the shipped headers, lives in NRT/KRT.)

QUIRK — POLL_SEM (0xB3) accelerates the spin on TOP_SP. ctrl_poll_sem.h documents a "specialized semi-custom operation to help Neuron accelerate polling semaphores on TOP_SP": read up to 16 consecutive uint32_ts at addr, fold an ALU op (currently required to be Min, ctrl_poll_sem.h:86-88) with the first value as seed, and if nonzero write the result back to all 16 plus a GP register (n_read/result/op/addr/res_writeback_addr fields, ctrl_poll_sem.h:43-53; 64 B, static-asserted). It too carries the same events@4 block. The lowered SYNC spin-wait can run through POLL_SEM on TOP_SP. (POLL_SEM role OBSERVED in its header; its use in SYNC lowering INFERRED / MED.)

NOTE — the older native event primitive. The Tonga-generation (V1) ISA had a dedicated EventWait op (aws_tonga_isa_tpb_event_wait.h: wait_event_mode

• event_idx, validity-checked, 64 B) and an EventSet op (aws_tonga_isa_tpb_event_set.h: "sets a local TPB event"). The modern EVENT_SEMAPHORE (0xA0) generalizes both into one wait+set nop carrying the events block — which is exactly why SYNC lowers onto it. (HIGH / OBSERVED.)

So the "device handler" for SYNC_BARRIER is the EVENT_SEMAPHORE (0xA0) / POLL_SEM (0xB3) HW path over the tpb_events array @0x0 — there is no opcode-0xD5 dispatch arm.

6 · The three-barrier boundary (within-core vs cross-core vs DMA)

All three barriers are 64-B pseudo-ops carrying the common header + the 8-B events block; they differ in barrier-specific payload and scope. Field inventories compile-verified (all sizeof == 64):

6.1 · The byte-exact SYNC vs CORE diff

This is the distinction the page exists to confirm. From the two headers + probe:

reserved0(52) − reserved1(44) = 8 == CORE.id(4) + CORE.semaphore(4) (probe prints the delta). SYNC_BARRIER is exactly CORE_BARRIER with the {id, semaphore} u32 pair removed and folded into the reserved padding — same header, same events block, at the same offsets. The removed pair is the counted-semaphore handle that SYNC does not need: it counts nothing, it sets and polls per-engine event flags. (HIGH / OBSERVED — both headers read, compile-verified.)

6.2 · The DMA barrier is orthogonal

PSEUDO_DMABARRIER (0xC3) is not a rendezvous barrier — it is a memory-consistency fence. The header doc (aws_neuron_isa_tpb_pseudo_dma_barrier.h:15-28) explains: TPB SBUF reads/writes are staggered across partitions, so a WR_DONE event means only partition 0 finished; a DMA could read a high partition before partition 0 and break the staggering. DmaBarrier before a DmaTrigger prevents that race. KRT lowers it to Nop(NUM_ROWS) — not to EVENT_SEMAPHORE. (The two NOP opcodes are ENGINE_NOP = 0x9f (common.h:238) and NOP = 0xa4 (common.h:244).) (HIGH / OBSERVED.)

6.3 · The scope ladder

0xD5 SYNC_BARRIER : ENGINES of ONE core           (finest;  EVENTS, no id/sema)
0xD8 CORE_BARRIER : PCORES of one VNC             (cross-core; SEMAPHORE + id)
NCFW 4-step barrier: RANKS/DIES of a collective   (coarsest; counted sema + dma_engines_bitmap)
0xC3 DMA_BARRIER  : (orthogonal) SBUF write-ordering Nop fence (no rendezvous)

(opcodes + doc comments + field inventories HIGH / OBSERVED; the relative ordering engines < cores < ranks is the architectural reading over the verbatim doc text, INFERRED-STRONG / MED.) The cross-die SBUF→SBUF path the coarser barriers gate is documented in RDMA Cross-Die SBUF→SBUF P2P.

7 · Use-case, liveness, and timeout

Use-case — "loop-on-chip flow" (verbatim): at a loop boundary inside a single NeuronCore, all engines that ran the loop body (PE produced, ACT/POOL/DVE consumed, SP orchestrated) must reach a common point before the next iteration / before a cross-core or DMA step. SYNC_BARRIER is the compiler's inline marker for that intra-core all-engine join — the fine-grain partner of CORE_BARRIER:

• align the engines of one core (loop iteration, before a DMA) → SYNC (0xD5)
• align the pcores of a VNC (collective phase, fused LNC group) → CORE (0xD8)
• order SBUF writes before a DMA read (memory consistency) → DMA (0xC3)

(the "loop-on-chip" use is the verbatim doc; the engine-roles narrative is the architectural reading over the NEURON_ENGINE enum — INFERRED / MED.)

GOTCHA — a missed event blocks forever. The SYNC_BARRIER struct has no timeout, no error-code, no max-spin field (only header/events/reserved0, §1), and there is no sync_barrier validator (§1). The event WAIT_MODE variants (0x6/0x7/0xe/0xf) have no "wait-or-timeout" form — an event wait is an unconditional spin until the bit reaches its target state. Consequence: if a participating engine never EVT_SETs its bit, the waiters block indefinitely. Liveness is a compiler/runtime guarantee (every engine emits its matching EVT_SET), not enforced by the instruction. The coarser cross-rank barrier carries timeouts at the host layer — but that is the cross-rank scope, not the on-chip SYNC pseudo-op. (absence of timeout in struct + in every event WAIT_MODE HIGH / OBSERVED; the indefinite-spin consequence INFERRED / MED.)

8 · Reimplementation checklist

For a Vision-Q7-compatible GPSIMD engine, SYNC_BARRIER is never decoded on device — it is a host/compiler artifact. A rebuild must:

1. Emit the 64-B pseudo at every intra-core, all-engine loop boundary: header.opcode = 0xD5, the events block (one seed wait/set pair), reserved0[52] zeroed. No id, no semaphore — that distinguishes it from 0xD8.
2. Allocate one tpb_events index (0..255, the 0x0 window) per participating engine of the core (PE/ACT/POOL/DVE/TPB_SP[/TOP_SP]).
3. Lower to a per-engine set of EVENT_SEMAPHORE (0xA0, CTRL_ES) ops: each engine EVT_SET_COMPLETEs (update_mode 0x11) its own bit and WAIT_FOR_EVT_SET (wait_mode 0x6, optionally _THEN_CLEAR 0x7) on the peers' bits — all updates on _Complete (es_updates_on_completion). Optionally route the TOP_SP spin through POLL_SEM (0xB3).
4. Drop the 0xD5 opcode entirely from the device dispatch table — the microcode never sees it.
5. Guarantee liveness in the rewriter: every participating engine must emit its matching EVT_SET, or the unconditional event spin deadlocks.

See also

• PSEUDO_CORE_BARRIER (0xD8) — the cross-core (pcores of a VNC) sibling; counted 32-bit semaphore + id. The byte-exact diff target of §6.1.
• PSEUDO_DMABARRIER (0xC3) — the orthogonal SBUF staggered-write fence, lowered to Nop(NUM_ROWS) by KRT.
• TriggerCollective (0xC8) — the canonical collective pseudo-op that carries the same events block.
• TriggerCollective2 + Ext (0xD9/0xDA) — the SBUF-tensor v2 collective carrier.
• Collective-Type + cc_op Enum Reference — the full COLLECTIVE_TYPE/ALU_OP/CCE_OP tables.
• The Unified Collective-Communication Architecture — how the barriers fit the collective/SPMD whole.
• RDMA Cross-Die SBUF→SBUF P2P — the EVT/SEM windows (read@0x1000/set@0x1400/inc@0x1800/dec@0x1C00, 256 entries, 4-B stride) at the cross-die scope.
• The Multicore API (8-core SPMD) · The 8-Core SPMD Execution Model + Teardown — the multicore context the within-core barrier operates inside.