test(rccl): add host microtests for memory management by aduforest · Pull Request #13609 · ROCm/rocm-systems

aduforest · 2026-10-15T23:21:23Z

JIRA ID: AICOMRCCL-2798

What

Host-only microtests for memory management: mem_manager.cc (teardown, Suspend/Resume with the peer handle exchange, the public entry points), allocator.cc (ncclMemAlloc/ncclMemFree, ncclSpace, ncclShadowPool), include/alloc.h (host/device/VMM alloc and free, side streams, the allocation tracker) and mnnvl.cc (ncclMnnvlCheck). They mostly extend test/mem_manager/MemManagerTests.cpp (a few reject-path checks overlap).

How

One binary per unit: rccl-UnitTestsMicro{MemManager,Allocator,Mnnvl,Alloc}. Alloc, Allocator and Mnnvl share fakes/alloc_micro_fakes.cc (logging, cuMem, getBusId and the memory-manager calls); MemManager's bootstrap/proxy/group seams are in fakes/mem_manager_micro_fakes.cc. HIP goes through fakes/hip_fakes.cc, which gains hooks for device sync, stream-ordered pools, priority streams and hipDeviceGet (defaults unchanged); rocmwrap_fakes.cc gains an init counter.
Pins 15 production bugs as KnownBug_ tests (AICOMRCCL-3036..3046). In the ROCm 7.12 arm the ncclCuMemHostAlloc macro also expands to an ncclCuMemHostAllocDebug nothing defines (AICOMRCCL-3047).

Coverage

Lines: mem_manager.cc 877/915 (909/915 with rccl-HostUnitTests), allocator.cc 363/363, alloc.h 658/663, mnnvl.cc 39/39.

Mutation Testing

146 mutations on the hipified sources, 145 killed; the survivor is equivalent.

Test Plan

MemManager 52/52, Allocator 37/37, Mnnvl 8/8, Alloc 51/51 in the standalone build, each also clean × 20 shuffled.
Allocator also with ROCM_VERSION=60400, and Mnnvl/Alloc with HIP_FABRIC_API and the ROCm 7.12 names stood in locally.

Cover entry teardown, Suspend/Resume with the peer handle exchange and re-import, and the public Suspend/Resume/MemStats entry points in rccl-UnitTestsMicroMemManager. hip_fakes gains hooks for the device sync, stream-ordered pool and priority-stream calls. Co-Authored-By: Claude <noreply@anthropic.com>

Cover the ncclMemAlloc/ncclMemFree VMM and fallback paths, the ncclSpace offset allocator and the ncclShadowPool object pool in rccl-UnitTestsMicroAllocator. hip_fakes gains a hipDeviceGet hook. Co-Authored-By: Claude <noreply@anthropic.com>

Cover ncclMnnvlCheck's cuMem, fabric-handle and vPoD-state gates, clique and NVL-domain discovery, and (when HIP_FABRIC_API is defined) the FABRIC export probe in rccl-UnitTestsMicroMnnvl. Co-Authored-By: Claude <noreply@anthropic.com>

Cover the shutdown and skip-free gates, side streams, host and device calloc/realloc, the cuMem VMM alloc/map/free helpers, the address-range walk and the allocation tracker in rccl-UnitTestsMicroAlloc. Co-Authored-By: Claude <noreply@anthropic.com>

List the four new binaries in one order in run_host_tests.sh and the README, and drop unused locals from the tests. Co-Authored-By: Claude <noreply@anthropic.com>

aduforest · 2026-10-15T23:21:59Z

@argus review

therock-pr-bot · 2026-10-15T23:22:26Z

✅ All Policy Checks Passed

Check	Status	Details
📝 PR Description	✅ Pass	—
⛔ Forbidden Files	✅ Pass	—
🧪 Unit Test	✅ Pass	—
🚫 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?

The in-build micro targets inherit rccl's HIP_VMM_UNCACHED_MEMORY, which makes ncclCuMemAlloc, ncclMemAlloc and Resume request hipMemAllocationTypeUncached. Also include the standard headers the tests use. Co-Authored-By: Claude <noreply@anthropic.com>

Where HIP has the FABRIC handle type, hipify renames CU_MEM_HANDLE_TYPE_FABRIC and CU_DEVICE_ATTRIBUTE_HANDLE_TYPE_FABRIC_SUPPORTED and p2p.h no longer defines them, so the tests resolve whichever name the build has. Co-Authored-By: Claude <noreply@anthropic.com>

aduforest · 2026-10-16T02:39:03Z

@argus review

aduforest · 2026-10-16T02:41:25Z

@argus review

amd-argus

ARGUS

Findings

5 more finding(s), not posted inline

projects/rccl/test/host/alloc-test.cc:903 (Low) CuMemFree_FailuresStillFreeTheRange pins the size-0 cuMemAddressFree bug without a BUG marker
projects/rccl/test/host/mem-manager-test.cc:73 (Low) Global operator delete replaced with std::free but operator new is not, so an ASan build reports alloc-dealloc-mismatch
projects/rccl/test/host/alloc-test.cc:466 (Low) CallocAndRealloc_ReportMallocFailure mallocs ~2^63 bytes, which ASan aborts on instead of returning NULL
projects/rccl/test/host/mem-manager-test.cc:914 (Low) Exchange_FabricEntryWithoutAnExportedHandleIsSentAnyway asserts on uninitialized malloc bytes
projects/rccl/test/host/allocator-test.cc:294 (Low) MemAlloc_/MemFree_ tests are not gated on ROCM_VERSION >= 70000 though allocator.cc's VMM path and rocmLibraryInit are

This PR changes 4,210 lines across 15 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: Looks approval ready.

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

hipMemLocationTypeHost is not declared before ROCm 7.12, which broke the ROCm 7.0.2 OCI build; spell it as its value. Also mark the size-0 cuMemAddressFree pin in ncclCuMemFree as a BUG, and let ASan return NULL for the oversized-malloc test. Co-Authored-By: Claude <noreply@anthropic.com>

…tests Gate the ncclMemAlloc/ncclMemFree VMM tests on ROCm 7.0 like allocator.cc and check the plain HIP path below it. Poison the mem_manager malloc seam so the unset FABRIC handle bytes are deterministic, and replace operator new alongside operator delete so ASan sees a matching pair. Co-Authored-By: Claude <noreply@anthropic.com>

aduforest · 2026-10-16T14:30:49Z

Addressed the five non-inline findings in 650b9d3 and 02a7531:

alloc-test.cc:903: CuMemFree_FailuresStillFreeTheRange now carries a BUG marker (same shared exit as ncclCuMemFreeAddr).
mem-manager-test.cc:73: operator new is replaced with a malloc-based one alongside operator delete.
alloc-test.cc:466: the binary sets allocator_may_return_null=1 through __asan_default_options, as init-test.cc does.
mem-manager-test.cc:914: the mem_manager malloc seam now poisons fresh blocks, and the test expects the poison byte (mutating the shareableHandleValid check is killed).
allocator-test.cc:294: the VMM ncclMemAlloc/ncclMemFree tests are under #if ROCM_VERSION >= 70000, with a plain hipMalloc/hipFree test below it (30/30 built at 60400).

Also fixed the ROCm 7.0.2 OCI build: hipMemLocationTypeHost is not declared there, so alloc-test.cc spells it as its value.

aduforest · 2026-10-16T14:30:54Z

@argus review

amd-argus

ARGUS

This PR changes 4,241 lines across 15 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.}

…microtests Keep both sides: develop's g_hipMemset hook next to the memory-manager hooks in hip_fakes.cc, and RomeModels/RmaProxy next to the four memory binaries in the registration files. Co-Authored-By: Claude <noreply@anthropic.com>

aduforest · 2026-10-16T16:40:17Z

@argus review

amd-argus

ARGUS

This PR changes 4,240 lines across 15 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.}

…easons, hip_fakes drift asserts - The 15 pinned bugs are <Fn>_KnownBug_<What> with one-line comments. - The empty-pool shadow lookup pin asserts SIGSEGV instead of any death. - ASSERT_HOOK_MATCHES_PROD for the 8 new hip_fakes hooks. - allocateSpilled fails loud; g_watchedDelete reset in TearDown. - The copied ncclHashPointer is pinned with rccl_assert_source_line. - CMake/README state the real own-binary reasons for MemManager and Mnnvl. Co-Authored-By: Claude <noreply@anthropic.com>

aduforest · 2026-10-16T21:19:21Z

@argus review

amd-argus

ARGUS

This PR changes 4,250 lines across 15 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 once the two ASan gaps are closed: DEATH_BY_SEGV at allocator-test.cc:840-841, and the nothrow operator new overloads in mem-manager-test.cc.

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

…rotests Co-Authored-By: Claude <noreply@anthropic.com>

aduforest · 2026-10-16T21:40:44Z

@argus re-review

amd-argus

ARGUS

This PR changes 4,270 lines across 15 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.}

…comrccl-2798-memory-host-microtests # Conflicts: # projects/rccl/test/host/CMakeLists.txt # projects/rccl/test/host/MICROTEST_README.md # projects/rccl/test/host/fakes/hip_fakes.cc # projects/rccl/test/host/run_host_tests.sh # projects/rccl/tools/scripts/test_runner/configs/ci-precheckin.json # projects/rccl/tools/scripts/test_runner/lib/test_executor.py

aduforest · 2026-10-19T18:04:06Z

@argus review

amd-argus

ARGUS

This PR changes 4,295 lines across 15 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.}

- The logging, cuMem, getBusId and memory-manager seams that alloc-test.cc, allocator-test.cc and mnnvl-test.cc each defined move to one fakes/alloc_micro_fakes.{h,cc}; the mem-manager binary's fakes move to fakes/mem_manager_micro_fakes.{h,cc}. Both are in the README map. - ncclOsGetPageSize comes from os_fakes.cc and rocmLibraryInit from rocmwrap_fakes.cc, which gains a call counter. - Hooks installed mid-test use ScopedHook; death tests restore death_test_style. - KnownBug_ pins name their Jira ticket. - test_executor.py lists only this PR's binaries. Co-Authored-By: Claude <noreply@anthropic.com>

…comrccl-2798-memory-host-microtests

aduforest · 2026-10-19T21:10:28Z

@argus review

amd-argus

ARGUS

This PR changes 4,357 lines across 21 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.}

dlamd1dai · 2026-10-19T23:45:20Z

+  }
+}
+
+TEST_F(MemManagerMicrotest, Suspend_OffloadBackupFailureStopsAtThatEntry) {


Suspend_OffloadBackupFailureStopsAtThatEntry doesn't lock the stop. Track prepends, so later is walked first and is the entry whose backup fails. Because the g_hipHostMalloc hook fails every call, the test still passes if Suspend records the error and continues instead of goto fail: offload's backup fails the same way, both entries stay Active, unmaps_ stays empty and released stays 0.

Could you track a scratch entry before the offload entry and assert it stays Active and is never unmapped? Failing only the first hipHostMalloc and asserting it ran once would also work.

dlamd1dai · 2026-10-19T23:45:21Z

+  std::free(dirty);
+}
+
+TEST_F(AllocMicrotest, Realloc_GrowsKeepingContentsAndZeroingTheTail) {


Realloc_GrowsKeepingContentsAndZeroingTheTail and IbMalloc_PageAlignedAndZeroed (line 417) only see dirty memory if the allocator hands back the block LeaveDirtyBlock freed. That holds with glibc today, and ASan fills new blocks with 0xbe. With a different allocator, the new block can come back zeroed and the tests would pass without the memset. A seam that returns a 0xff-filled buffer would remove that dependency. Optional.

dlamd1dai · 2026-10-19T23:45:23Z

+}  // namespace
+
+// utils.cc's hash; with hbits == 0 it returns bucket 0 instead of shifting by 64.
+uint64_t ncclHashPointer(int hbits, void* key) {


The CMake checks pin three lines of ncclHashPointer but not the seed, uintptr_t h = reinterpret_cast<uintptr_t>(key);. If upstream changes that line, configure still passes while the copy's bucket placement drifts from misc/utils.cc. Could you add it to the rccl_assert_source_line list?

The extra hbits == 0 branch is fine as far as I can tell. allocator.cc sets hbits = 4 before it ever hashes, and the empty-pool crash is the null table either way.

dlamd1dai · 2026-10-19T23:45:24Z

+TEST_F(AllocMicrotest, CuMemHostAlloc_MapsHostMemoryForTheDeviceAndTheCpu) {
+  void* p = nullptr;
+  hipMemGenericAllocationHandle_t h = nullptr;
+  ASSERT_EQ((ncclCuMemHostAlloc)(&p, &h, 100), ncclSuccess);


The parentheses in (ncclCuMemHostAlloc) skip the macro. On ROCm >= 7.12 that macro expands to ncclCuMemHostAllocDebug, which nothing defines (AICOMRCCL-3047). So a translation unit that uses the macro would still fail to link while this suite stays green. A compile-only pin, or a line saying the macro is intentionally left unpinned, would make that explicit.

aduforest and others added 5 commits October 15, 2026 16:17

test(rccl): add host microtests for src/mnnvl.cc

f34ce04

Cover ncclMnnvlCheck's cuMem, fabric-handle and vPoD-state gates, clique and NVL-domain discovery, and (when HIP_FABRIC_API is defined) the FABRIC export probe in rccl-UnitTestsMicroMnnvl. Co-Authored-By: Claude <noreply@anthropic.com>

test(rccl): order the memory-management microtest registrations

747d800

List the four new binaries in one order in run_host_tests.sh and the README, and drop unused locals from the tests. Co-Authored-By: Claude <noreply@anthropic.com>

github-actions Bot added the project: rccl label Oct 15, 2026

systems-assistant Bot added the organization: ROCm label Oct 15, 2026

amd-argus Bot reviewed Oct 16, 2026

View reviewed changes

aduforest and others added 2 commits October 16, 2026 07:29

amd-argus Bot reviewed Oct 16, 2026

View reviewed changes

aduforest marked this pull request as ready for review October 16, 2026 15:01

aduforest requested review from a team October 16, 2026 15:01

amd-argus Bot reviewed Oct 16, 2026

View reviewed changes

Comment thread projects/rccl/test/host/allocator-test.cc Outdated

Comment thread projects/rccl/test/host/mem-manager-test.cc Outdated

test(rccl): ASan-safe SEGV pin and nothrow operator new in memory mic…

c6f58d9

…rotests Co-Authored-By: Claude <noreply@anthropic.com>

amd-argus Bot reviewed Oct 16, 2026

View reviewed changes

amd-argus Bot reviewed Oct 19, 2026

View reviewed changes

aduforest and others added 2 commits October 19, 2026 14:10

Merge remote-tracking branch 'origin/develop' into users/aduforest/ai…

4f8b79f

…comrccl-2798-memory-host-microtests

amd-argus Bot reviewed Oct 19, 2026

View reviewed changes

dlamd1dai reviewed Oct 19, 2026

View reviewed changes

Conversation

aduforest commented Oct 15, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

What

How

Coverage

Mutation Testing

Test Plan

Uh oh!

aduforest commented Oct 15, 2026

Uh oh!

therock-pr-bot Bot commented Oct 15, 2026

✅ All Policy Checks Passed

Uh oh!

aduforest commented Oct 16, 2026

Uh oh!

aduforest commented Oct 16, 2026

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Findings

Uh oh!

aduforest commented Oct 16, 2026

Uh oh!

aduforest commented Oct 16, 2026

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Uh oh!

aduforest commented Oct 16, 2026

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Uh oh!

aduforest commented Oct 16, 2026

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

aduforest commented Oct 16, 2026

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Uh oh!

aduforest commented Oct 19, 2026

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Uh oh!

aduforest commented Oct 19, 2026

Uh oh!

amd-argus Bot left a comment

Choose a reason for hiding this comment

Uh oh!

dlamd1dai Oct 19, 2026

Choose a reason for hiding this comment

Uh oh!

dlamd1dai Oct 19, 2026

Choose a reason for hiding this comment

Uh oh!

dlamd1dai Oct 19, 2026

Choose a reason for hiding this comment

Uh oh!

dlamd1dai Oct 19, 2026

Choose a reason for hiding this comment

Uh oh!

This branch has not been deployed

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

aduforest commented Oct 15, 2026 •

edited

Loading