CAST NET/IB multi-segment GPUDirect RDMA by atulkulk · Pull Request #10889 · ROCm/rocm-systems

atulkulk · 2026-08-28T15:03:19Z

Motivation

CAST NET/IB (NCCL_NET=IB-CAST) had the same first-segment DMA-BUF limitation as classic NET/IB. Multi-segment buffers could not complete GPUDirect RDMA on the CAST plugin.

Technical Details

Private IbCastRegMrDmaBufMultiSeg; side table after the CAST CTS FIFO; 64-byte CTS and 32-byte AINIC inline CTS unchanged.
Segmented send splits the WRR-selected QP chunk at independent local and remote boundaries; per-segment IFlush fences every non-zero receive and every segment it overlaps.
Completion-size records distinguish the logical transfer size from the final physical slice.
Only one segmented group is admitted per QP slot, and failed posts return borrowed WR credits.
Connect cap NCCL_IB_CAP_MULTISEG is encoded in the fixed-size devName trailer; no connection-metadata field or CTS-size change.
net.cc does not branch on the CAST plugin. IbCastRegMrDmaBuf does the same fd < 0 range walk as classic. The segmented send and flush stay in p2p.cc.
The sender may register a multi-segment source without the peer cap. The receive side table is allocated only for a multi-segment peer.
Sender and flush QP depths grow only while NCCL_MULTI_SEGMENT_REGISTER is on. With the switch on, the sender depth is 769 work requests and the flush depth is 1024, matching classic. With the switch off, both stay at the develop depths.

Issue Tracking

JIRA ID: AICOMRCCL-1526

Test Plan

Reuse NetIbMultiSegment* under NCCL_NET=IB-CAST.
Single-segment CAST traffic must stay bit-compatible; peers without multiseg cap must stage.

Test Result

The focused host boundary/RMA helper suite passed 28/28.
The two-node CAST P2P wire suite passed 12/12 with NCCL_NET=IB-CAST, including per-segment registration, whole-buffer and cross-boundary transfers, independent offsets, completion size, multi-receive flush, maximum-segment handling, and single-segment regression.
These are completed hardware runs; cross-boundary cases were neither skipped nor declined to staging.

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.
CHANGELOG: deferred to a follow-up PR.

therock-pr-bot · 2026-08-28T15:05:45Z

✅ All Policy Checks Passed

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	⚠️ Warning	Error: Source/code files changed without an accompanying unit test. Expected: add at least one test file named like `test_<name>.py` / `test_<name>.cpp` (or `<name>_test.`). Current:* code file(s) changed: `projects/rccl/src/ce_coll.cc`, `projects/rccl/src/dev_runtime.cc`, `projects/rccl/src/include/alloc.h`, `projects/rccl/src/include/ce_fault_inject.h`, `projects/rccl/src/include/rocmwrap.h` (+32 more); no test file found
🚫 Draft PR	🔜 To Be Enabled	—
🚩 Feature Flag	🔜 To Be Enabled	—
📊 Code Coverage	🔜 To Be Enabled	—

🎉 All policy checks passed!

📖 Need help? See the Policy FAQ for details on every check and how to fix failures.

🙋 Wish to Override Policy?

Copilot

🟡 Changes recommended

Segmented CTS and flush traffic can exhaust undersized queues, failed posts leak remap entries, and the new test configuration leaves CAST flushing disabled.

Once you've addressed the issues Copilot identified, you can request another Copilot review.

Pull request overview

Adds multi-segment GPUDirect RDMA support to the CAST NET/IB transport.

Changes:

Adds CAST-specific multi-segment registration and capability negotiation.
Implements segmented sends, completion sizing, and multi-segment flushing.
Extends MPI coverage and CAST test-runner configuration.

File summaries

File	Description
`projects/rccl/tools/scripts/test_runner/configs/net_ib_transport.json`	Adds CAST multi-segment test configuration.
`projects/rccl/test/transport/NetIbMPI/NetIbMultiSegmentMPITests.cpp`	Runs registration through the CAST helper and verifies completion sizes.
`projects/rccl/test/transport/NetIbMPI/NetIbMultiSegmentHelpers.hpp`	Dispatches registration to the selected IB plugin.
`projects/rccl/src/transport/rccl_ib_multiseg.h`	Declares the CAST registration helper.
`projects/rccl/src/transport/net.cc`	Enables CAST multi-segment proxy registration.
`projects/rccl/src/transport/net_ib_cast/reg.cc`	Implements segmented MR registration and cleanup.
`projects/rccl/src/transport/net_ib_cast/p2p.cc`	Implements segmented transfers, side-table publication, completion sizing, and flushing.
`projects/rccl/src/transport/net_ib_cast/p2p_cast.h`	Adds remote segmented-layout detection.
`projects/rccl/src/transport/net_ib_cast/connect.cc`	Negotiates capabilities and allocates/registers supporting resources.
`projects/rccl/src/transport/net_ib_cast/connect_cast.h`	Enforces connection metadata compatibility.
`projects/rccl/src/transport/net_ib_cast/common_cast.h`	Adds segmented wire structures, MR state, and WR sizing.

Review details

Files reviewed: 11/11 changed files
Comments generated: 5
Review effort level: Balanced

💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.

amd-argus

ARGUS

Findings

3 finding(s) that fall outside the diff and could not be anchored to a line

projects/rccl/src/transport/net_ib_cast/p2p.cc:839 (Medium) IbCastIsend has no isendFail recovery, so a send-build error leaves the slot occupied and records no fatal error, unlike the classic sibling
projects/rccl/test/transport/NetIbMPI/NetIbMultiSegmentMPITests.cpp:366 (Low) Both flush tests make the flush assertion conditional, and MultiSegmentFlushTouchesEverySegment never verifies the payload
projects/rccl/tools/scripts/test_runner/configs/net_ib_transport.json:1232 (Low) The ib-cast to IB-CAST rename is a real coverage fix, but 97 lowercase entries remain in 7 other config files and the descriptions here are stale

Bottom line: Worth settling before approval: WR_IMM_SEGMENTED_FLAG is tested on the BY_ORDER path at p2p.cc:1506, where imm_data is the raw request id, so a receive can report 0 bytes once fifoHead passes 0x400000.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus · 2026-09-14T16:55:55Z

 // The alignment for IB writes that is required to make LL and LL128 protocols work
 #define IB_WRITE_CHUNK_ALIGNMENT 128

+static ncclResult_t IbCastMultiSendSegmented(struct ncclIbSendComm* comm, int slot, int nqps, int startQpIndex,


IbCastMultiSendSegmented duplicates about 26 lines of IbCastMultiSend byte for byte: the ENABLE_FAULT_INJECTION block at p2p.cc:349-359 against p2p.cc:607-617, and the trailing scheduler-update and log block at p2p.cc:373-387 against p2p.cc:661-675. Those two extract into static helpers cleanly and would stop the copies drifting. Roughly 40 further lines are near-identical but would need parameters.

Agreed it is worth extracting, but I'd keep it out of this PR to avoid touching IbCastMultiSend's develop body; it is on the follow-up list with the flush re-walk above.

amd-argus

ARGUS

Findings

1 finding(s) that could not be posted as inline comments

projects/rccl/test/transport/NetIbMPI/NetIbMultiSegmentMPITests.cpp:366 (Low) All three flush tests hide a disabled flush behind if (freq != nullptr), so they pass whether or not the new per-segment WRs ran MultiSegmentFlushSelectsSegmentMr wraps its completion check in if (freq != nullptr) at NetIbMultiSegmentMPITests.cpp:366, as do line 408 and, as flushReq, line 458. IbCastIflush returns early without a request when flushEnabled is 0 (p2p.cc:1180), so the new per-segment flush WRs can go entirely unexercised while all three tests pass. cast_multisegment now forces NCCL_GDR_FLUSH_DISABLE=0, but flushEnabled still depends on GDR at runtime. Could a null request become an explicit skip with a reason instead?

Bottom line: Worth settling before approval: the WR_IMM_SEGMENTED_FLAG test at p2p.cc:1506 runs in the BY_ORDER handler, where imm_data is a raw request id, so a receive reports 0 bytes once fifoHead passes 0x400000. Unchanged since the last round.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus · 2026-09-16T19:50:02Z

+  // Side table immediately after ctsFifo so one covering MR registers both.
+  // Receiver RDMA-writes this only when nSegments>1 and the peer advertised
+  // NCCL_IB_CAP_MULTISEG.
+  struct ncclIbSegLayout segLayoutFifo[NET_IB_MAX_REQUESTS][NCCL_NET_IB_MAX_RECVS];


segLayoutFifo adds NET_IB_MAX_REQUESTS * NCCL_NET_IB_MAX_RECVS * 704 bytes to every ncclIbSendComm, and ncclIbRemSegLayout at common_cast.h:737 the same to every ncclIbRecvComm, unconditionally.

segLayoutFifo 256 * 8 * 704 B = 1.375 MiB per send comm remSegLayout 256 * 8 * 704 B = 1.375 MiB per recv comm covering MR ctsFifo + segLayoutFifo ~= 1.5 MB pinned per device (connect.cc:1133)

Comms are per channel per peer, so this multiplies in both directions even when multi-segment is never enabled. Could the table be allocated on the first nSegments > 1 registration instead?

The receive side is fixed in e276df6265, matching classic: ncclIbRecvComm::remSegLayout is now a pointer allocated (and its per-device MR registered) only when the peer advertised NCCL_IB_CAP_MULTISEG, and the irecv fill returns early without it. The sender's segLayoutFifo stays inline because the peer RDMA-writes it at ctsFifo + sizeof(ctsFifo) through the CTS covering MR, so moving it would need a second address and rkey in the connect metadata; the covering MR only grows to include it when the switch is on.

amd-argus

ARGUS

Findings

1 finding(s) that could not be posted as inline comments

projects/rccl/src/transport/net_ib_cast/common_cast.h:708 (Low) The %32 assert cannot fail under alignas(64), and the classic sibling's absolute-size assert was dropped sizeof(struct ncclIbSegLayout) % 32 cannot fail: the struct is alignas(64) at common_cast.h:410, so its size is always a multiple of 64. The classic copy pairs it with an absolute size assert (net_ib/common.h:613), which is the one that carries weight here, since 704 is baked into sgeSide.length at p2p.cc:1025 and the slot stride at p2p.cc:1022.

This PR changes 1,203 lines across 13 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.

Bottom line: Worth settling before approval: the segmented send path reports an over-long receive size under BY_ORDER (p2p.cc:1684), which is the default scheme on AINIC, where this PR's own cast_multisegment config runs.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus

ARGUS

This PR changes 1,204 lines across 13 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.

Bottom line: Worth settling before approval: four High findings from earlier rounds are still open, led by the rkey read from unvalidated peer-writable memory at p2p.cc:277 and the unencoded QP-sharing commId at p2p.cc:306.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus

ARGUS

This PR changes 1,553 lines across 13 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.

Bottom line: LGTM.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus

ARGUS

This PR changes 1,570 lines across 13 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.

Bottom line: Nothing new worth raising; still open from earlier: p2p.cc:717. Worth clearing before approval.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

atulkulk · 2026-10-21T20:31:08Z

@argus re-review

amd-argus

ARGUS

This PR changes 1,628 lines across 14 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.

Bottom line: Worth settling before approval: the NetIbMultiSegmentMPITest.* filters at net_ib_transport.json:753 and 785 also run OldPeer_SingleRkeyCrossBoundaryTransfer, whose rank-divergent ASSERT_EQ desyncs the two ranks.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

amd-argus · 2026-10-21T20:53:17Z

+        {
+          "name": "Cast_MultiSegment_All",
+          "description": "CAST registration, asymmetric/cross-boundary transfer, multi-receive, flush, limits, and single-segment regressions",
+          "test_filter": "NetIbMultiSegmentMPITest.*"


High

NetIbMultiSegmentMPITest.* here and at line 753 also runs OldPeer_SingleRkeyCrossBoundaryTransfer. ncclParamMultiSegmentRegister() is cached at 1 by then (this entry sets it and earlier tests connect), so rank 1 fails the ASSERT_EQ at NetIbMultiSegmentMPITests.cpp:461 while rank 0 continues into SyncSkip and the connect collectives. Could both filters append -*OldPeer*, with a CAST twin of the standalone entry classic has at line 574?

Resolved by ed63d70fd2.

amd-argus

ARGUS

This PR changes 1,645 lines across 20 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.

Bottom line: LGTM.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

BertanDogancay · 2026-10-22T18:44:24Z

+      for (struct ibv_send_wr* w = comm->wrs; w != NULL && w != bad_wr; w = w->next) posted++;
+      // RCCL: Isend armed one event per QP. A prefix must set sentData or isendFail drops it and frees the request.
+      if (posted > 0) {
+        for (int r = 0; r < nreqs; r++) reqs[r]->send.sentData[qpIndex] = true;


If this post fails after the NIC took the earlier writes those writes are not signaled. Setting sentData here keeps the completion isend already counted so the test waits for the one that never comes. The flush path above frees the request in this case. Can we do the same?

- net-ib-cast: add multi-segment P2P wire option Per-segment DMA-BUF registration and segmented CAST send/recv for uniform GPU multi-segment buffers. Dispatch netIbRegMrMultiSeg to IbCast on NCCL_NET=IB-CAST. Route NetIb MPI test registration through the matching plugin helper under IB-CAST. - net-ib-cast: harden multi-segment P2P correctness Preserve mixed-version wire compatibility and make segmented completion sizing, flush coverage, queue credits, and failure cleanup correct under multi-receive traffic.

…lushes - net-ib-cast: free remap on post failure and enable GDR flush Release BY_INDEX remap entries when post_send or fault inject fails, and set NCCL_GDR_FLUSH_DISABLE=0 on the CAST multi-segment suite. - net-ib-cast: share CTS layout checks and ncclCalloc for multi-seg MRs Use the classic CTS validator and allocate the multi-seg wrapper with ncclCalloc, matching the single-segment path. - net-ib-cast: wait for the CTS side table and keep posted irecv requests - net-ib-cast: bound CTS side-table wait and drain failed sends - net-ib-cast: drain a failed flush the way isend does IbCastIflush used NCCLCHECK on post_send, so a failed post leaked the request slot. Add iflushFail, keep FAILED when a prefix posted, and return the post error when nothing was accepted. - test: require a flush request when GDR flush is enabled A null iflush request used to skip the wait, so a disabled flush still passed the multi-segment flush tests. - test: require at-cap success before over-cap reject ExceedsMaxSegmentsRejected passed if every register returned ncclInvalidUsage. Register NCCL_IB_MAX_SEGMENTS first so a policy decline cannot masquerade as the cap. - net-ib-cast: drop unused ibCastMrForRange and dead arExtra needSizesWr is always true on the segmented send, so arExtra never changed extraImmWr. Move multiseg.h into the include group. - test: fail CAST Generic when NET_REG_COMPLETE is missing Skip would stay green if IbCastRegMrDmaBufMultiSeg were reverted. Also zero each CAST flush WR as it is filled. - cast: record telemetry on segmented send - cast: size flush QP for multi-segment recv A segmented irecv can post one flush WR per overlapping segment, which exceeds NET_IB_MAX_REQUESTS (256) on the loopback flush QP.

…safe modes - net-ib-cast: pick the zero-length rkey by offset like classic - net-ib-cast: pin devName's connect-metadata offset with a static_assert - test: cover CAST multi-segment receives with optional recv completion - net-ib-cast: decline unsafe multi-seg modes and fail stuck requests Decline multi-segment registration under BY_ORDER and QP sharing, and only advertise the cap when MULTI_SEGMENT_REGISTER is on. Release events for unposted send QPs, and return errors from irecv CTS and partial flush failures instead of spinning in Test. Bound the peer's CTS nreqs; verify hinted recvs after the barrier. - net-ib-cast: honour the multi-segment switch on the use side; cover =0, BY_ORDER and QP sharing Plugin: - Mask the peer's MULTISEG cap with the local NCCL_MULTI_SEGMENT_REGISTER switch on both connect and accept, so a =0 sender never waits on a side table and a =0 receiver never writes one. - Register the receiver's side-table MR only when MULTISEG was negotiated. - Size the sender, CTS and flush QP send queues for multi-segment only when the switch is on; =0 keeps the previous depths, including on AINIC recovery rebuilds. - The segmented builder uses its validated copies of the remote segment starts and rkeys instead of rereading the live side table. - Drop two irecv checks the pre-validation already covers, restore the flush INFO line to its previous form (the TRACE above logs the WR count), and wrap an over-long line. - Make the multi-segment registration header comment plugin-neutral. Tests: - MultiSegmentRegisterOff_DeclinesRegistration routes through the plugin-matching registration entry, so it builds and runs under IB-CAST. - New CastUnsafeMode_DeclinesRegistration expects ncclInvalidUsage under BY_ORDER matching or QP sharing. - UBR_MultiSegment.Generic keys the CAST arm on the comm's loaded plugin (so the ROCM-IB alias and the unset default count) and skips the CAST arm when the dma-buf export gate is compiled out. - The CAST flush tests also require NCCL_GDR_FLUSH_DISABLE=0. - The at-cap check agrees across ranks before the over-cap call, and the send side no longer repeats the receive-size check. Runner (net_ib_transport.json): - Keep the existing CAST jobs' NCCL_NET values and cast_base as before; only the multi-segment jobs set NCCL_NET=IB-CAST. - New cast_multisegment_register_off (=0) and cast_multisegment_unsafe_modes (BY_ORDER, QP sharing) jobs. They extend cast_base with the multi-segment env, so they run only their own tests. - net-ib-cast: move multi-segment P2P and registration into their own files Mirror the classic layout: p2p_multiseg_cast.h holds the multi-segment send, CTS side-table, irecv and flush paths behind small hooks in p2p.cc, and rccl_reg_multiseg_cast.cc holds multi-segment MR registration and deregistration. p2p.cc, reg.cc and p2p_cast.h shrink back toward develop. Register the receiver's side-table staging MR with local write only, drop an unreachable transport branch, and restore unrelated develop comments. Accept, PD teardown and post-failure cleanup move to the next commit. - net-ib-cast: harden accept, PD teardown and post-failure cleanup Take the PD reference only after allocation succeeds, and make IbCastDestroyBase safe on a partially initialised base. Route every IbCastAccept failure through fail and close the stage comm there. Free the send remap entry when a post or injected fault fails, and clear the request and record a fatal error on the irecv error path. - net-ib-cast: register one MR per segment when regMrDmaBuf is called with fd < 0 The CAST plugin walks the range itself, matching classic IB, so generic registration does not branch on the plugin name. - net-ib-cast: keep the multi-segment P2P bodies in p2p.cc Mirror the classic layout: a separate header was a second file to carry on every upstream sync and was included in the middle of p2p.cc. - net-ib-cast: let the sender register a multi-segment source without the peer cap The side table describes the receiver; a sender splits a local multi-segment source onto the receiver's single rkey, as classic does. CTS offload, BY_ORDER and QP sharing still decline on both sides. - net-ib-cast: allocate the recv side table only for a multi-segment peer and bound the segmented nreqs reload The receiver's side-table source is allocated and registered only when the peer advertised the cap, as classic does. The segmented send rejects a peer-written nreqs below one, and the CAST multi-segment configs drop the RCCL_IB_P2P_DISABLE_CTS setting that does not apply to them.

… post cannot complete - net: size the CAST flush QP for a chain on every proxy step CAST used the same one-extra-chain depth as classic. The proxy can leave one direct-GDR flush posted on every step. - net: keep a CAST multi-segment request after a prefix post A failed post that already accepted work requests was freed before Test() could drain it. - net: free a CAST flush request when the signaled read was not posted A partial multi-read flush post leaves only unsignaled work, so arming an event makes Test wait for a completion that never arrives. A host fixture fails the signaled tail and checks that the request is freed. The CAST multi-segment runner runs the old-peer case in its own process, and every wildcard that would cache the switch first excludes it. - net: free a CAST segmented send when the signaled tail was not posted A partial segmented post set sentData, so isendFail kept an event for a completion that never comes. sentData stays clear, and isendFail drops that QP's event and frees the request, as iflush does. Host tests fail the tail on either QP.

amd-argus

ARGUS

This PR changes 1,780 lines across 20 files. A change this size is hard to review
well, for a person and for this bot alike: reviewers lose the thread, defects hide in the
volume, and the feedback arrives too late to be cheap to act on. Splitting it into a
stacked pull request
gives each piece a reviewable diff and keeps the history. This review covers the whole change
regardless.

Bottom line: Nothing new worth raising; still open from earlier: common_cast.h:746. Worth clearing before approval.

_{Comment @argus review to re-run, or @argus re-review after pushing.}

github-actions Bot added the project: rccl label Aug 28, 2026

systems-assistant Bot added the organization: ROCm label Aug 28, 2026

atulkulk changed the base branch from develop to users/atulkulk/mseg/07-hybrid-stress-ci August 28, 2026 15:09

atulkulk force-pushed the users/atulkulk/mseg/07-hybrid-stress-ci branch from 9cbe456 to b22240c Compare August 28, 2026 15:47

atulkulk force-pushed the users/atulkulk/mseg/08-cast-netib-multisegment branch 2 times, most recently from db4af3b to fe92cd0 Compare August 28, 2026 16:08

atulkulk force-pushed the users/atulkulk/mseg/07-hybrid-stress-ci branch from 0357bb8 to 652c77d Compare September 7, 2026 09:30

atulkulk force-pushed the users/atulkulk/mseg/08-cast-netib-multisegment branch from fe92cd0 to 8273b40 Compare September 7, 2026 09:31

atulkulk requested a balanced review from Copilot September 7, 2026 19:01

Copilot started reviewing on behalf of atulkulk September 7, 2026 19:02

View session

Copilot AI reviewed Sep 7, 2026

View reviewed changes

atulkulk force-pushed the users/atulkulk/mseg/07-hybrid-stress-ci branch from 652c77d to ad6a208 Compare September 14, 2026 15:13

atulkulk force-pushed the users/atulkulk/mseg/08-cast-netib-multisegment branch from 8273b40 to b0ecd11 Compare September 14, 2026 15:13

atulkulk marked this pull request as ready for review September 14, 2026 15:18

atulkulk requested review from a team as code owners September 14, 2026 15:18

amd-argus Bot reviewed Sep 14, 2026

View reviewed changes

BertanDogancay force-pushed the users/atulkulk/mseg/07-hybrid-stress-ci branch from ad6a208 to fa2939d Compare September 16, 2026 18:39

BertanDogancay force-pushed the users/atulkulk/mseg/08-cast-netib-multisegment branch from b0ecd11 to 186893f Compare September 16, 2026 18:40

amd-argus Bot reviewed Sep 16, 2026

View reviewed changes

atulkulk force-pushed the users/atulkulk/mseg/07-hybrid-stress-ci branch from fa2939d to 37bbee8 Compare September 23, 2026 12:09

atulkulk force-pushed the users/atulkulk/mseg/08-cast-netib-multisegment branch from 186893f to 5f6cc15 Compare September 23, 2026 12:10

amd-argus Bot reviewed Sep 23, 2026

View reviewed changes

atulkulk force-pushed the users/atulkulk/mseg/07-hybrid-stress-ci branch from 37bbee8 to c75b007 Compare October 1, 2026 21:15

atulkulk force-pushed the users/atulkulk/mseg/08-cast-netib-multisegment branch from 5f6cc15 to cebf36b Compare October 1, 2026 21:16

amd-argus Bot reviewed Oct 1, 2026

View reviewed changes

Comment thread projects/rccl/src/transport/net_ib_cast/p2p.cc Outdated

Comment thread projects/rccl/src/transport/net_ib_cast/p2p.cc Outdated

Comment thread projects/rccl/src/transport/net_ib_cast/p2p.cc Outdated

atulkulk force-pushed the users/atulkulk/mseg/08-cast-netib-multisegment branch from cebf36b to a783a22 Compare October 2, 2026 01:31

atulkulk force-pushed the users/atulkulk/mseg/07-hybrid-stress-ci branch from c75b007 to 650ffc7 Compare October 2, 2026 01:31

atulkulk force-pushed the users/atulkulk/mseg/08-cast-netib-multisegment branch from a783a22 to b4b6177 Compare October 2, 2026 09:14

atulkulk force-pushed the users/atulkulk/mseg/07-hybrid-stress-ci branch from 5e17680 to d656db2 Compare October 16, 2026 02:31

atulkulk force-pushed the users/atulkulk/mseg/08-cast-netib-multisegment branch from e276df6 to 534cecc Compare October 16, 2026 02:44

atulkulk force-pushed the users/atulkulk/mseg/07-hybrid-stress-ci branch from d656db2 to 19cfc1a Compare October 18, 2026 01:05

atulkulk force-pushed the users/atulkulk/mseg/08-cast-netib-multisegment branch from 534cecc to 1b821ec Compare October 18, 2026 01:05

amd-argus Bot reviewed Oct 18, 2026

View reviewed changes

atulkulk force-pushed the users/atulkulk/mseg/07-hybrid-stress-ci branch from 19cfc1a to fb0f790 Compare October 19, 2026 22:14

atulkulk force-pushed the users/atulkulk/mseg/08-cast-netib-multisegment branch from 1b821ec to dd9c9cb Compare October 19, 2026 22:14

atulkulk force-pushed the users/atulkulk/mseg/07-hybrid-stress-ci branch from fb0f790 to f5999ad Compare October 20, 2026 04:21

atulkulk force-pushed the users/atulkulk/mseg/08-cast-netib-multisegment branch from dd9c9cb to 95db53f Compare October 20, 2026 04:22

atulkulk force-pushed the users/atulkulk/mseg/07-hybrid-stress-ci branch from f5999ad to 63e0055 Compare October 20, 2026 05:02

atulkulk force-pushed the users/atulkulk/mseg/08-cast-netib-multisegment branch from 95db53f to defdad6 Compare October 20, 2026 05:02

atulkulk force-pushed the users/atulkulk/mseg/07-hybrid-stress-ci branch from 63e0055 to 77a0bad Compare October 21, 2026 16:07

atulkulk force-pushed the users/atulkulk/mseg/08-cast-netib-multisegment branch from defdad6 to a1d86ba Compare October 21, 2026 16:08

amd-argus Bot reviewed Oct 21, 2026

View reviewed changes

atulkulk force-pushed the users/atulkulk/mseg/08-cast-netib-multisegment branch from 2dd6c58 to ed63d70 Compare October 21, 2026 21:31

amd-argus Bot reviewed Oct 21, 2026

View reviewed changes

atulkulk force-pushed the users/atulkulk/mseg/07-hybrid-stress-ci branch from 77a0bad to 226e5c3 Compare October 22, 2026 13:28

atulkulk force-pushed the users/atulkulk/mseg/08-cast-netib-multisegment branch from ed63d70 to 6280387 Compare October 22, 2026 13:28

atulkulk force-pushed the users/atulkulk/mseg/07-hybrid-stress-ci branch from 226e5c3 to 5de3006 Compare October 22, 2026 17:32

atulkulk force-pushed the users/atulkulk/mseg/08-cast-netib-multisegment branch from 6280387 to 2fffae1 Compare October 22, 2026 17:33

BertanDogancay reviewed Oct 22, 2026

View reviewed changes

atulkulk added 4 commits October 23, 2026 13:05

atulkulk force-pushed the users/atulkulk/mseg/08-cast-netib-multisegment branch from 2fffae1 to c3cc8f6 Compare October 23, 2026 16:35

atulkulk force-pushed the users/atulkulk/mseg/07-hybrid-stress-ci branch from 5de3006 to 66629ab Compare October 23, 2026 16:35

amd-argus Bot reviewed Oct 23, 2026

View reviewed changes

Conversation

atulkulk commented Aug 28, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Motivation

Technical Details

Issue Tracking

Test Plan

Test Result

Submission Checklist

Uh oh!

therock-pr-bot Bot commented Aug 28, 2026

✅ All Policy Checks Passed

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

🟡 Changes recommended

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Findings

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

amd-argus Bot Sep 14, 2026

Choose a reason for hiding this comment

Uh oh!

atulkulk Oct 16, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Findings

Uh oh!

Uh oh!

Uh oh!

amd-argus Bot Sep 16, 2026

Choose a reason for hiding this comment

Uh oh!

atulkulk Oct 16, 2026

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Findings

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Uh oh!

amd-argus Bot left a comment

atulkulk commented Aug 28, 2026 •

edited

Loading