Repository navigation
[AICOMRCCL-1851] [RCCL] test: cover rma_proxy cpuAccessSignals MR type - #12756
Conversation
✅ All Policy Checks Passed
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
ed9f738 to
0f66325
Compare
There was a problem hiding this comment.
Copilot review overview
🟡 Changes recommended
Graph allocations leak during fixture teardown and between test iterations, accumulating during repeated runs and failing leak detection.
Review effort: Balanced
Findings: 1
Open (1)
What changed in this PR
Adds host-only regression coverage for RCCL’s existing RMA-proxy memory registration fix, without changing production code.
Changes:
- Tests host and GDRCopy registration types, strict plugin rejection, and ordering flags.
- Adds a resettable synchronous
hipMemsetfake. - Integrates the dedicated test binary into builds, execution, categories, and documentation.
| File | Description |
|---|---|
| projects/rccl/test/test_categories_micro_rma_proxy.yaml | Registers test categories. |
| projects/rccl/test/host/run_host_tests.sh | Runs the new binary. |
| projects/rccl/test/host/rma-proxy-test.cc | Adds five regression tests. |
| projects/rccl/test/host/MICROTEST_README.md | Documents the binary and usage. |
| projects/rccl/test/host/fakes/hip_fakes.h | Declares the memset hook. |
| projects/rccl/test/host/fakes/hip_fakes.cc | Implements and resets the hook. |
| projects/rccl/test/host/CMakeLists.txt | Defines the dedicated host-test target. |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
![amd-argus[bot]](https://avatars.githubusercontent.com/in/4740020?s=60&v=4){kind=small}
|
@mch @pvallem — review request for this one. @dlamd1dai tagging you for visibility. Adds host-only regression coverage for the Six tests in a new Evidence it would have caught the defect: with Review comments from Copilot and argus-bot are addressed and resolved. |
There was a problem hiding this comment.
ARGUS
Bottom line: Worth settling before approval: whether ncclRmaProxyCtxAllocGraph nulls what it frees on its error path, since the new reject-all test is the first to run FreeGraph() against a partially built ctx_.
Comment @argus review to re-run, or @argus re-review after pushing.
There was a problem hiding this comment.
ARGUS
Findings
1 finding(s) that could not be posted as inline comments
projects/rccl/src/rma/rma_proxy.cc:233(Low) ncclCalloc is passed a byte count where it expects an element count, over-allocating cpuAccessSignalsHost 8xncclCalloctakes an element count, but rma_proxy.cc:233 (and :164) passsignalsBufSize, which is alreadysignalSlots * sizeof(uint64_t). Both targets areuint64_t*, so each allocates 8x: 512 bytes for a 64-byte buffer at nRanks=8, numRmaSig=1. Over-allocation only, and thefree()stays correct. rma_ce.cc:82 passessignalSlotsfor the same buffer. Outside this diff, so noting it as a follow-up rather than something to fix here.
Bottom line: Looks approval ready.
Comment @argus review to re-run, or @argus re-review after pushing.
|
Thanks — re-review on On the Agreed it is a follow-up and not for this PR: it is over-allocation only (no overflow, and |
|
Note on the Its true state is
Neither log contains a compiler diagnostic, a link error or any test output — the build never got far enough to emit one. This is not specific to this PR. The last ten runs of For what it is worth, this PR adds host-only test sources plus a CMakeLists entry and a test-categories yaml — nothing that the perf gate measures. |
78a696b to
0bdd4eb
Compare
|
Rebased onto current Two conflicts, both additive list entries resolved by keeping both sides:
No test-source changes were needed. The fix under test survives the 2.32 sync unchanged at Revalidated on a compute node against the rebased tree: refreshed the hipified sources ( |
|
The three checks that went red after the rebase ( Root cause. coco's console log — which is uploaded as an artifact rather than printed in the job log, which is why the job log shows only a bare
Evidence that it is not the rebase. The invocation is byte-identical to the last green run on this branch; the only thing that changed is the resolved coco ref:
And it is repo-wide — every RCCL PR whose coco cluster jobs actually ran today after ~15:24Z fails both clusters the same way, e.g. runs 37643636626 (PR #13030), 37643848856, 37645593861, 37648852799. I confirmed the artifact for 37643636626 carries the identical error string. Nothing to fix on this branch, and rerunning will not help until either the workflow's The host microtest itself was revalidated against the rebased tree on a compute node (refreshed hipify, Debug build, 3 repeats, |
mch
left a comment
mch
left a comment
There was a problem hiding this comment.
Can you tidy up the cmake and remove the assertion of the fake behvaiour?
|
Removed assertions on fake behaviour:
Kept, deliberately:
CMake: source-list preamble collapsed to one line, inline comments dropped from the target block, now matching the Revalidated on a compute node, Debug, 3 repeats: 6/6 pass. Negative control with Your review carried no inline anchors, so if anything else reads as fixture-asserting-fixture, point at it and I will take it out. |
|
Correction to my previous comment: your review did carry inline anchors — eight of them — and I had read a stale enumeration when I said otherwise. All eight are now addressed and resolved, with a per-thread reply. Summary across 5eb14f3 and 5115f8d:
Validation on a compute node, Debug (what CI builds), 3 repeats: 6/6 pass. Negative control with @mch ready for another look. |
|
The
Two things settle it:
The underlying failure is a job-step launch fault, not an RCCL one: The test binary never starts — the PMI/daemon fan-out dies before Not rerunning, for the same reason as before: a rerun would land on the same cluster state and cost cycles without new information. Happy to rerun if a CI owner wants it once the OCI nodes are healthy. |
|
Follow-up on the other two red checks, now that
Both ran the complete size sweep to 2 GiB with Same shape as the OCI failure reported above: errors confined to the 4-node topology, with the develop baseline affected identically. No
This PR adds test-only files under |
5115f8d to
053d6bc
Compare
|
@mch — rebased onto current Rebase. The branch was 257 commits behind. Two conflicts, both the same additive-list shape as the last rebase, both resolved by keeping both sides:
No test changes were needed. The fix still sits unchanged at Revalidated on a compute node (job 71037,
Review state. All eight of your inline anchors were addressed across One thing I did not touch, to keep the diff tight: line 799 of |
053d6bc to
bc276c0
Compare
Add host-only microtests for the MR-registration-type fix in ncclRmaProxyCtxAllocGraph (src/rma/rma_proxy.cc). When GDRCopy is disabled -- the default on ROCm, since NCCL_GDRCOPY_ENABLE defaults to 0 -- allocMemCPUAccessible hands back host-NUMA pinned memory rather than device memory. The pre-fix code unconditionally registered cpuAccessSignalsDev as NCCL_PTR_CUDA, so the net plugin rejected the MR. Upstream fixed this in NCCL 2.30.7 (NVIDIA/nccl PR #2187) by deriving the type from the GDR handle: non-NULL means GDRCopy-mapped device memory (NCCL_PTR_CUDA), NULL means host memory (NCCL_PTR_HOST). The line was previously uncovered, and because the host path is the default path on AMD, a regression would silently break every RMA-proxy run without GDRCopy. The tests pivot the real allocMemCPUAccessible on the extern global ncclGdrCopy rather than mocking the allocator, so the GDR handle the production ternary branches on is the one production actually produced. A FakeRma vtable records every (addr, size, type, mrFlags) registration; in the strict-plugin test it rejects NCCL_PTR_CUDA for any pointer that the hipHostMalloc seam really returned, so the failure mode mirrors the defect instead of restating the assertion. Five cases: both GDR arms of the cpuAccessSignals registration, a strict plugin that refuses a mistyped MR, flushBufDev staying NCCL_PTR_CUDA on both arms, and NCCL_NET_MR_FLAG_FORCE_SO preserved on both arms. This lands as its own binary, rccl-UnitTestsMicroRmaProxy, rather than joining rccl-UnitTestsMicro: that target already resolves ncclRmaProxyConnectOnce, ncclRmaProxyRegister and ncclRmaProxyDeregister from fakes/dev_runtime_micro_fakes.cc, and rma_proxy.cc defines all three. Also upgrades the fail-loud hipMemset stub in fakes/hip_fakes.cc to a g_hipMemset seam whose default really memsets, so a buffer production zeroes reads back as zero. No existing host test referenced hipMemset.
…ture teardown
RmaProxyAllocGraphTest only destructed comm_->memPermanent, which releases the
persistent/in-progress queues but none of the heap allocations
ncclRmaProxyCtxAllocGraph makes: cpuAccessSignals (allocMemCPUAccessible),
cpuAccessSignalsHost (ncclCalloc) and flushBufDev (hipExtMallocWithFlags).
Deleting ctx_ does not free its raw-pointer members. The two arm-looping tests
additionally overwrote those pointers on their second AllocGraph() call.
Add FreeGraph(), which undoes the allocations in reverse order, each step
null-guarded so a partially completed AllocGraph is still cleaned up. It
deregisters the recorded MRs and pairs the real allocator with the real
deallocator -- freeMemCPUAccessible, selected by the same GDR handle the unit
under test branches on -- then drops the freed pointer from the hipHostMalloc
provenance set so a recycled address cannot make the next arm's "not host
memory" check pass or fail for the wrong reason. Called from TearDown (so it
also runs after an early return from a failed assertion) and at the top of each
loop iteration in FlushBufferStaysCudaOnBothGdrArms and
ForceStrongOrderingPreservedOnBothGdrArms.
ncclRmaProxyDestroyContext is deliberately not used: it ends in free(rmaProxyCtx),
which is incompatible with the fixture's unique_ptr ownership.
The GDR arm's free path reaches ncclCudaFree, which queries the allocation range
before freeing, so SetUp now opens the g_hipMemGetAddressRange seam alongside the
other GDRCopy-on gates.
Verified on a compute node, Debug (what run_host_tests.sh builds):
- 5/5 tests, --gtest_repeat=3, rc=0
- -fsanitize=address with ASAN_OPTIONS=detect_leaks=1, --gtest_repeat=3:
5/5, rc=0, no LeakSanitizer report
- negative control, same ASan build of the previous revision:
"SUMMARY: AddressSanitizer: 397397 byte(s) leaked in 24 allocation(s)"
…review nits Review feedback on the rma_proxy cpuAccessSignals MR-type microtests. Add PluginRefusesEveryRegistration_AllocGraphFails. The three tests that assert ncclRmaProxyCtxAllocGraph returns ncclSuccess are only evidence if a refused registration can still make it fail; nothing demonstrated that the NCCLCHECK around the cpuAccessSignals registration propagates. A rejectAll flag on FakeRma refuses every registration and the new test asserts AllocGraph does not succeed. Assert the registered size on the GDRCopy-on arm too. ncclGdrCudaCalloc rounds its allocation up to GPU_PAGE_SIZE, so that is the one arm where a registration sized from the padded block instead of signalsBufSize would go unnoticed. Correct the props.ptrSupport comment. It claimed the NCCL_PTR_DMABUF bit shows the registrations are steered by memory type, but ncclRmaProxyRegMrSym tests ncclParamDmaBufEnable() first and the fixture pins that to 0, so the bit is never read. Note the same gate on the EXPECT_FALSE(viaDmaBuf) assertion, which checks the gate held rather than any property of the fix. Also assert the host-memory provenance in the strict-plugin test and name the assertion in ForceStrongOrderingPreservedOnBothGdrArms. Verified on a compute node, Debug, --gtest_repeat=3: 6/6, rc=0; and under -fsanitize=address with ASAN_OPTIONS=detect_leaks=1: 6/6, rc=0, no LeakSanitizer report.
…to a helper Review follow-ups on the rma_proxy cpuAccessSignals MR-type microtest: - FlushBufferStaysCudaOnBothGdrArms and ForceStrongOrderingPreservedOnBothGdrArms carried a byte-identical set-gdr / FreeGraph / AllocGraph preamble. Moved it into the fixture as ReAllocGraphForArm, called under ASSERT_NO_FATAL_FAILURE so a failure inside the helper still stops the test. - The destroyContext vtable slot was installed but never checked. TearDown now asserts destroyCtxCalls == 0, which is the property the function's own comment states: ncclRmaProxyCtxAllocGraph leaves failure cleanup to the caller's ncclRmaProxyDestroyContext and must not destroy the context itself. Validated on a compute node, Debug and ASan/LSan, 3 repeats each: 6/6, rc=0, no sanitizer report.
… cmake
Review feedback: the microtest asserted on its own test doubles in several
places, which tests the fixture rather than the unit under test.
Removed:
- the EXPECT_EQ on the handle tag inside FakeRma::Rma, and with it the
now-vestigial Handle::Kind tag;
- the hostAddrs_ provenance checks in four tests -- the
cpuAccessSignalsGdrHandle premise checks already pin which allocator arm
ran;
- EXPECT_TRUE(rmaNet_.regs.empty()) in the reject-everything test, which
only restated how the fake was configured;
- EXPECT_FALSE(reg->viaDmaBuf), self-described as "not a property of the
fix", and with it the unread viaDmaBuf plumbing;
- the EXPECT_ wrappers around the cleanup calls in FreeGraph().
hostAddrs_ itself stays: it is the input to strictMemType, the mechanism
behind StrictPluginRejectsWrongMemoryType_AllocGraphStillSucceeds.
EXPECT_EQ(0, rmaNet_.destroyCtxCalls) in TearDown stays: it asserts
production behaviour, that AllocGraph must not tear the context down.
CMakeLists: collapsed the source-list preamble to one line and dropped the
inline comments on the target block, matching the rccl-UnitTestsMicroDiagnostics
house style.
Revalidated on a compute node, Debug: 6/6 pass; with rma_proxy.cc:229 reverted
to the pre-fix NCCL_PTR_CUDA literal, exactly GdrCopyOff_RegistersCpuAccess
SignalsAsHost and StrictPluginRejectsWrongMemoryType_AllocGraphStillSucceeds
fail and the other 4 pass -- the regression anchors survive the trim.
…cycle Review feedback, inline threads on the previous head. Lifecycle: the fixture no longer hand-builds an ncclRmaProxyCtx and calls the file-static ncclRmaProxyCtxAllocGraph directly, with a hand-written FreeGraph() mirroring ncclRmaProxyDestroyContext. It now goes through the entry points a production client uses -- ncclRmaProxyCreateContext (which runs CtxAlloc then CtxAllocGraph) and ncclRmaProxyDestroyContext -- so production owns the context and its cleanup, and the cleanup no longer has to be kept in step by hand. The vtable gains a createContext trampoline and the comm carries rmaState.rmaProxyState.ncclRma. PluginRefusesEveryRegistration_AllocGraphFails becomes PluginRefusesCpuAccessSignalsRegistration_CreateContextFails: rejectAll is replaced by rejectAfterNReg, set to 1, which lets the signalsDev registration from CtxAlloc through and refuses the cpuAccessSignals one specifically, so the negative control still pins the NCCLCHECK under test rather than an earlier one. CMakeLists: trimmed the source list to what the link needs -- env_fakes, comm_fakes, mem_manager_fakes, recorder_fakes, strongstream_stubs, bootstrap_stubs and archinfo.cc are gone -- and replaced the per-entry annotations with the one-line pointer to the README map that the Init and Enqueue lists use. Comments: dropped the enumeration of rma_proxy.cc's params (already drifted), replaced the file:line citations with file-and-symbol, deleted the trailing explanations on three includes, and reduced the own-binary rationale to a pointer at the MICROTEST_README entry that carries it in full. Validated on a compute node: Debug, 3 repeats, 6/6 pass. With rma_proxy.cc:229 reverted to the pre-fix NCCL_PTR_CUDA literal, exactly GdrCopyOff_RegistersCpuAccessSignalsAsHost and StrictPluginRejectsWrongMemoryType_AllocGraphStillSucceeds fail. Clean under ASan with detect_leaks=1.
bc276c0 to
f335478
Compare
…proxy-test Merging develop brought in #12756, which defines a local ncclSetThreadName stub in rma-proxy-test.cc. This branch had already moved that symbol onto its owning fakes file (fakes/nccl_fakes.cc), which rccl-UnitTestsMicroRmaProxy also links, so the two definitions collided at link time: ld.lld: error: duplicate symbol: ncclSetThreadName(std::thread&, char const*, ...) Remove the local stub and note in the surrounding comment that symbols with an owning fakes file must not be redefined in this TU. Co-authored-by: Pi Agent
TheRock Submodule Bump ActivityNewest first
|
Motivation
ncclRmaProxyCtxAllocGraphregisterscpuAccessSignalsDevwith the net plugin.That buffer comes from
allocMemCPUAccessible, which returns host-NUMA pinnedmemory whenever GDRCopy is off — and
NCCL_GDRCOPY_ENABLEdefaults to0, soon ROCm the host path is the default path, not an edge case.
Before the fix the registration type was hardcoded
NCCL_PTR_CUDA, so the netplugin rejected the MR. Upstream fixed this in NCCL 2.30.7 (NVIDIA/nccl PR #2187)
by deriving the type from the GDR handle. RCCL
developis at NCCL 2.31.2 andalready carries the fix, so no production code changes here — the line was
simply uncovered, and a regression would silently break every RMA-proxy run
without GDRCopy. RCCL has already had to patch two neighbouring symptoms of the
same "these signals may be host memory" confusion (
rcclParamRmaProxyUseDMABUFand the
props_tmpDMA-BUF strip in the same function), which is why pinningthis behaviour down is worth doing.
This PR adds the regression coverage.
Technical Details
New host-only microtest
test/host/rma-proxy-test.ccplus its own binary,rccl-UnitTestsMicroRmaProxy. Host-only: it links gtest + fmt, notlibrccl.so,libamdhip64.soorlibhsa-runtime64.so, and needs no GPU.It gets a dedicated binary rather than joining
rccl-UnitTestsMicrobecausethat target already resolves
ncclRmaProxyConnectOnce,ncclRmaProxyRegisterand
ncclRmaProxyDeregisterfromfakes/dev_runtime_micro_fakes.cc, andrma_proxy.ccdefines all three — linking the unit under test there is aduplicate symbol.
MICROTEST_README.mdsanctions one binary per unit, andrccl-UnitTestsMicroInit/...Enqueueare the precedent.ncclRmaProxyCtxAllocGraphis file-static, so the test TU#includes thehipified
.ccvia a build-timeRMA_PROXY_CC_PATHmacro — the documented microidiom, same as the sibling
rma-proxy-progress-test.cc.Two seams make this a test of the fix rather than a restatement of it:
allocMemCPUAccessibleis astatictemplate ingdrwrap.h, so it cannot bestubbed at link time. That is a feature here: the test runs the real helper
and pivots it on the real discriminator, the extern global
ncclGdrCopy. TheGDR handle the production ternary branches on is the one production actually
produced, not a mocked premise.
StrictPluginRejectsWrongMemoryType_AllocGraphStillSucceeds,FakeRmarejects
NCCL_PTR_CUDAfor any pointer present in a set populated by thehipHostMallocseam. "Is this host memory?" is answered by real provenance,so the test models the defect rather than asserting the conclusion.
Also in this PR:
fakes/hip_fakes.{h,cc}:hipMemsetwas a fail-loud stub with no hook, whichwould have aborted this function's
cudaMemset(flushBufDev, ...)underCUDACHECK. Upgraded to ag_hipMemsetseam whose default really memsets. Noexisting host test referenced
hipMemset.run_host_tests.sh: added the new binary to the run list. Without this the CIhost-test job would build it and never run it, and still go green.
test/test_categories_micro_rma_proxy.yaml: category registration, structurallyidentical to
test_categories_micro_enqueue.yaml.Issue Tracking
JIRA ID: AICOMRCCL-1851
Upstream: NVIDIA/nccl PR #2187 (NCCL 2.30.7)
Test Plan
All compiles and all binary runs were done on compute nodes, never a login node.
test/hostproject) and in-tree.rma_proxy.cc— the exact TU the binary
#includes — back to the pre-PR-Rock: hip-tests installation path to remain same for linux/windows #2187 unconditionalNCCL_PTR_CUDA, rebuild, rerun, then restore and re-verify green.llvm-cov --show-branches=count: both arms of thecpuAccessSignalsGdrHandle != NULLternary must be non-zero. Line hits aloneare not evidence — the bug lives entirely in which arm is taken.
--gtest_shuffle(CI shuffles), and a JUnit reportactually written.
Test Result
Green, in-tree and standalone: 5/5,
rc=0, andlddconfirms nolibrccl/libamdhip/libhsa.Pre-fix revert —
rc=1, 2 failed / 3 passed:Restoring the line returns 5/5. The second failure is the defect stated verbatim:
a strict plugin refusing a mistyped MR takes
ncclRmaProxyCtxAllocGraphtoncclInternalError.The three passes are as important as the failures — they show the suite is not
simply "everything red when the line changes". The GDRCopy-on arm, the flush
buffer and the
FORCE_SOflag are unaffected, so an over-correction to alwaysNCCL_PTR_HOSTwould be caught too.Branch coverage:
Shuffle: 5/5 and a valid
<testsuites>report at seeds 0, 1, 7, 42, 1851.Scope limit, stated plainly
This is host-microtest evidence. The real-world claim — that a net plugin stops
rejecting the MR — is modelled by a
FakeRmavtable, not exercised against alive IB plugin on multi-node hardware. For a one-line memory-type fix on a path
that is already the ROCm default, I judged a host microtest the proportionate
gate; a reviewer who wants an MPI/IB functional test should say so.
Submission Checklist