Add accumulate method for PortChannel - #784
Open
Changho Hwang (chhwang) wants to merge 9 commits into
Open
Conversation
Changho Hwang (chhwang)
force-pushed
the
chhwang/new-atomic-add
branch
from
August 4, 2026 17:20
4b21685 to
de4997b
Compare
Changho Hwang (chhwang)
changed the title
atomicAdd method for PortChannelaccumulate method for PortChannel
Changho Hwang (chhwang)
marked this pull request as ready for review
August 4, 2026 23:36
|
Azure Pipelines: There may be pipelines that require an authorized user to comment /azp run to run. |
Changho Hwang (chhwang)
force-pushed
the
chhwang/new-atomic-add
branch
3 times, most recently
from
August 6, 2026 08:22
ccfecde to
36ce485
Compare
Changho Hwang (chhwang)
changed the base branch from
main
to
chhwang/enhanced-fifo
August 7, 2026 19:07
Changho Hwang (chhwang)
force-pushed
the
chhwang/new-atomic-add
branch
from
August 7, 2026 19:07
36ce485 to
776cf75
Compare
Binyang Li (Binyang2014)
force-pushed
the
chhwang/enhanced-fifo
branch
from
August 12, 2026 23:07
e00cff7 to
bab7f8c
Compare
Binyang Li (Binyang2014)
force-pushed
the
chhwang/new-atomic-add
branch
from
August 12, 2026 23:07
776cf75 to
03fa28b
Compare
Binyang Li (Binyang2014)
force-pushed
the
chhwang/new-atomic-add
branch
from
August 12, 2026 23:10
03fa28b to
97f82d8
Compare
`TriggerData`, `TriggerFlag`, and `TriggerSync` were combinable bits, but the
device API only ever produced five of the seven combinations, and the proxy
tested them with three independent masks. Replace them with one opcode per
operation the API can express:
TriggerNone = 0 TriggerPut = 1 TriggerSignal = 2
TriggerFlush = 3 TriggerPutWithSignal = 4 TriggerPutWithSignalAndFlush = 5
Seven values fit the existing 3-bit field with one spare. `handleTrigger`
becomes a switch that names one operation per case and warns on anything
unknown, so a combination nothing emits is now unrepresentable, and a trigger
whose type field is unset is no longer a valid operation.
The old names are removed rather than redefined. Code that composed them with
`|`, or tested them with `&`, must fail to compile: `TriggerFlush` is 3, which
has the old `TriggerData` bit set, so a mask test would silently misbehave.
No behavior change. mp_unit_tests 29/29 and unit_tests 34/34 on H200.
Add `Connection::accumulate()` and `PortChannel::accumulate()`, which add a 64-bit signed value to remote memory. Unlike `updateAndSync()`, the caller needs to know only its own contribution, not the destination's current value. Carried by the `TriggerAccumulate` opcode, with the operand in the trigger's first word. FIFO readiness now comes from the lap-parity commit bit, so a zero operand traverses the proxy like any other trigger; transport tests interleave zero and nonzero accumulates to cover that path. The addition has to be a real read-modify-write at the destination, which not every transport can do: - IB: any number of writers, via RDMA fetch-and-add. Rejected in no-atomic mode, where the device has no RDMA atomics. - Ethernet: any number of remote writers. The receiving process performs the update, and its connections are serialized against each other; without that a 7-writer fan-in loses ~74% of updates. - CudaIpc on ROCm: any number of writers. The proxy runs a kernel on the connection's stream, which also orders it against a following signal. A kernel there runs even while the caller's kernel spins. - CudaIpc on CUDA: not supported, throws InvalidUsage. The host cannot read-modify-write device memory, and a proxy-launched kernel either cannot be scheduled while the caller's kernel spins (same context) or costs 2391 us per operation (separate context), against 19 us for a plain remote store. Tested on H200 (CUDA 12.9) and MI300X (ROCm 7.2), including an 8-rank fan-in test with 7 concurrent writers into one address.
Changho Hwang (chhwang)
force-pushed
the
chhwang/new-atomic-add
branch
from
August 14, 2026 13:59
97f82d8 to
ccab6d7
Compare
…' into chhwang/new-atomic-add # Conflicts: # src/core/connection.cc
Changho Hwang (chhwang)
changed the base branch from
main
to
fix/ethernet-receiver-close-race
August 20, 2026 08:04
| constexpr TriggerType TriggerFlush = 4; // Flush the connection. | ||
| constexpr TriggerType TriggerAccumulate = 6; // Add a value to remote memory. | ||
| constexpr TriggerType TriggerPutWithSignalAndFlush = 7; // Transfer data, signal, then flush. | ||
| // 5 is unassigned. |
Contributor
There was a problem hiding this comment.
Do we need to reserve 5 here? You mean user may build customized algo with old API?
Contributor
Author
There was a problem hiding this comment.
It's just unused so unassigned, but probably we'd better reserve for a future operation. Will do
| // not wait for the caller. | ||
| accumulateU64Kernel<<<1, 1, 0, *stream_>>>(dst, value); | ||
| MSCCLPP_CUDATHROW(cudaGetLastError()); | ||
| dirty_ = true; |
Contributor
There was a problem hiding this comment.
Why need to set dirty here? It will be reset some where?
|
|
||
| // Every peer terminates its socket in this process, so several recv threads can be here | ||
| // at once for one address. The read-modify-write below is not atomic, so serialize it | ||
| // against the other recv threads. |
Contributor
There was a problem hiding this comment.
Why there will be several threads here. For each device we just launch one recv thread, no?
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
2 participants
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.
Summary
Adds
Connection::accumulate()andPortChannel::accumulate()for remote signed 64-bit addition. Callers provide only their contribution, allowing concurrent writers to share a counter.The operation uses the
TriggerAccumulateFIFO opcode. Its operand occupies the trigger's first word. FIFO readiness comes from lap-parity commit bits, so zero-valued operands are valid.Transport support
Validation and ordering
Tests
Coverage includes:
Follow-up validation for the current head:
The original implementation was also validated on H200/CUDA 12.9 and MI300X/ROCm 7.2.