Repository navigation
Conversation
ariscript
requested review from
BKP,
artulab,
drprajap,
efric,
mawad-amd,
neoblizz and
nirvedhmeshram
as code owners
October 29, 2026 19:04
This was referenced Oct 29, 2026
Contributor
There was a problem hiding this comment.
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 Pythonrange(...)in… · New In Triton,range(...)must have a compile-time constant bound; `tl.load(self.view +… · Newtl.constexpris meant to be used as a parameter annotation in Triton kernel signatures, not… · NewH_SENDS_TOis 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
…iggered plans Gates become one int32 symmetric tensor: a counter and an arrival counter per chunk, then a completion counter per peer. - Chains poll the chunk's counter against writers x round, so mark_ready is a single release atomic: no last-writer election and no ready word. plan.counters is contiguous int32, so a kernel that signals per tile (a tritonBLAS epilogue) can increment it directly. - Every round, a rank's chains to each peer end with an ATOMIC ADD32 on the sending rank's completion counter for that peer. wait() then reads only local words; the progress kernel's remote reads and armed table are gone. TriggeredView gains sent() / wait_sent(). - Arrival signals are 32-bit (ADD32, 0x0f). release_kernel raises the counter to writers x round with atomic_max, so close() opens partially marked gates. - Round limit: writers x round <= 2^31 - 1. Tests pass at 2 and 4 ranks on MI300X (torch allocator). Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
The device view keeps the peers a rank sends to as bits of one int64 word; at 64 ranks building it failed with an int64 overflow from torch. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…nment Assigning a Python int to an element of the CUDA handle tensor copies from the host and blocks until the stream drains, so start() (and close()'s release) waited for all queued GPU work. fill_ on a one-element view is queued like any kernel. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
ariscript
force-pushed
the
arprakas/plan-gate-counters
branch
from
October 30, 2026 22:43
53d8704 to
8accc21
Compare
This branch has not been deployed
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.


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
counter per peer.
writers × round, somark_readyis a singlerelease atomic. There is no last-writer election and no separate ready word.
plan.countersis contiguous int32, so a kernel that signals per tile can increment it directlyinstead of calling
TriggeredView.mark_ready. That includes a tritonBLASEpilogueConfig(counter_buffer=...)(Add SignalView and WaitView tritonBLAS#81).
counter for that peer:
plan.wait()now reads only local memory; the progress kernel's remote reads are gone.TriggeredView.sent()/wait_sent(), for example beforeoverwriting their source.
writers × roundmust stay at or below 2³¹ − 1 (over 7M rounds at 304 writers).Testing (MI300X, torch allocator)
tests/unittests/test_plan.pyat 2 and 4 ranks. New tests:plan.counterswith a plain gpu-scopeatomic_addafterdebug_barrier, as the tritonBLAS epilogue does: 8 writers on one 4 MiB chunk, 10 rounds,no stale data;
wait_sentbefore reusing the source;