[ariscript]

First of three stacked follow-ups to #573, in the order agreed on that PR: gate layout (this PR),
then rounds (#617), then destinations and arrival groups (#618). They address @dsidler's review points
1 and 4.

What changes

• Gates are one int32 symmetric tensor: a counter and an arrival counter per chunk, then a completion
  counter per peer.
• The copy engine polls a chunk's counter against writers × round, so mark_ready is a single
  release atomic. There is no last-writer election and no separate ready word.
• plan.counters is contiguous int32, so a kernel that signals per tile can increment it directly
  instead of calling TriggeredView.mark_ready, for example a tritonBLAS EpilogueConfig(counter_buffer=...)
  (Add SignalView and WaitView tritonBLAS#81). That works only when the kernel's counter numbering matches the schedule (only
  some map_type / schedule-builder pairs do) and its barrier waits for stores (Triton with
  [Backend] Add a new TritionGPU barrier triton-lang/triton#9025); the docs spell out both.
• Every round, each peer queue's chains end with one copy-engine add on the sending rank's completion
  counter for that peer:
  • plan.wait() now reads only local memory; the progress kernel's remote reads are gone.
  • Kernels can wait inline with the new TriggeredView.sent() / wait_sent(), for example before
    overwriting their source.
  • The peers a rank sends to are one flag each in the handle tensor, so there is no rank limit.
  • Each peer's completion signal is queued separately, even after a failed arm, and the original error
    is the one raised. A peer whose signal cannot be queued is no longer waited for, and the plan's
    memory is kept (with a warning) rather than freed while its copies may still run.
• wait() also waits until this rank's own chunks have all their marks, so it never checks the mark
  count before the producer has finished.
• start() sets the round with a fill_ queued on the current stream. It used to assign an element of
  the handle tensor, which copies from the host and waited for all queued GPU work. Kernels that take
  the handle must run on that stream, or wait for it.
• Limit: writers × round must stay at or below 2³¹ − 1 (over 7M rounds at 304 writers).

Testing (MI300X, torch allocator)

• tests/unittests/test_plan.py at 2 and 4 ranks. New tests:
  • a kernel that increments plan.counters with a plain gpu-scope atomic_add after
    debug_barrier, as the tritonBLAS epilogue does: 8 writers on one 4 MiB chunk, 10 rounds,
    no stale data;
  • wait_sent before reusing the source, from a second kernel and from inside a persistent producer;
  • a completion signal that cannot be queued, alone or after a failed arm;
  • the round limit.
• View tests (Triton and Gluon) on one GPU, including the wait kernel waiting for this rank's marks and
  sent() past 64 ranks.
• Docs build.

[Copilot]

Warning

Copilot couldn't run its full agentic review because it didn't start before the timeout. Make sure your repository has a runner available, or add a copilot-code-review.yml file specifying one with the runs-on attribute. See the docs for more details.

Copilot review overview

4 open issues

• Same issue as the Triton path: the loop bound is loaded at runtime, but the Python range(...) in… · New
• In Triton, range(...) must have a compile-time constant bound; `tl.load(self.view +… · New
• tl.constexpr is meant to be used as a parameter annotation in Triton kernel signatures, not… · New
• H_SENDS_TO is stored in a single int64 header word (bitmask), which cannot represent `world_size… · New

What changed in this PR

Updates triggered-plan gate layout and device/host synchronization to use contiguous int32 counters (per-chunk marks + arrivals, and per-peer completion), enabling single-atomic mark_ready and local-only wait() progress checks.

Changes:

• Replace per-chunk “ready word + counter” gate rows with int32 counters/arrivals and per-peer completion counters (ABI v2).
• Add device APIs TriggeredView.sent() / wait_sent() and host-side completion signaling per peer queue.
• Update unit tests and docs to reflect the new gate tensor layout and semantics.

File	Description
iris/​experimental/​plan/​triton.py	Switches device ABI/layout to int32 gates and adds sent()/wait_sent() logic.
iris/​experimental/​plan/​gluon.py	Mirrors Triton TriggeredView changes for Gluon device code.
iris/​experimental/​plan/​triggered.py	Allocates/slices new gate tensor, adds round limit check, and switches wait logic to local-only gates.
iris/​experimental/​plan/​kernels.py	Replaces progress kernel with local-only wait kernel; updates release semantics.
iris/​experimental/​plan/​armer.py	Arms chains polling counter thresholds and signals int32 arrivals + per-peer completion.
tests/​unittests/​test_plan.py	Adds coverage for external counter increments, wait_sent, and round limit; updates arrivals assertions.
tests/​unittests/​test_plan_view.py	Updates view tests for new counter/arrival layout and adds sent()/wait_sent() coverage.
tests/​unittests/​test_plan_view_gluon.py	Updates Gluon view tests similarly and adds sent()/wait_sent() coverage.
docs/​conceptual/​triggered-plans.md	Updates conceptual model and memory-ordering text for counter-based gating and completion counters.
docs/​reference/​host/​plan.md	Exposes new device methods sent/wait_sent in API docs.
iris/​experimental/​plan/​__init__.py	Updates module docstring to mention new device APIs.

_{Review effort: Lite}

[nirvedhmeshram]

Thanks, left some comments, also if some of them will be addressed in #617 or #618, please feel free to ignore them here.

[nirvedhmeshram]

LGTM