From 85c83d5c1745648a4964f76a3f7d6e20500d6c80 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Tue, 1 Sep 2026 18:17:11 -0400 Subject: [PATCH 01/17] fix(rccl): link window mem onto memHead before GIN register (AICOMRCCL-2241) Anvil SDMA resolves user VA via ncclDevrGetLsaSelfAddr during regMrSym, which only walks memHead. Linking after GIN/RMA register hid in-flight windows (legal after ncclDevCommCreate). Unlink and deregister on the failure path. Co-authored-by: Cursor --- projects/rccl/src/dev_runtime.cc | 25 +++++++++++++++++++++---- 1 file changed, 21 insertions(+), 4 deletions(-) diff --git a/projects/rccl/src/dev_runtime.cc b/projects/rccl/src/dev_runtime.cc index 0cd9f6e4eec3..cc50a9c36ab7 100644 --- a/projects/rccl/src/dev_runtime.cc +++ b/projects/rccl/src/dev_runtime.cc @@ -851,6 +851,13 @@ static ncclResult_t symMemoryObtain(struct ncclComm* comm, CUmemGenericAllocatio } } + // Link before GIN/RMA registration. Plugins such as Anvil SDMA resolve the + // user VA through ncclDevrGetLsaSelfAddr, which only searches memHead (and + // the LSA flat range). Registering after ncclDevCommCreate is legal, so the + // in-flight mem must already be on the list. Unlink on the failure path. + mem->next = devr->memHead; + devr->memHead = mem; + if (devr->ginEnabled) { NCCLCHECKGOTO(symMemoryRegisterGin(comm, mem), ret, fail_mem_space_teams); } else { @@ -866,15 +873,24 @@ static ncclResult_t symMemoryObtain(struct ncclComm* comm, CUmemGenericAllocatio NCCLCHECKGOTO(symMemoryRegisterRma(comm, mem), ret, fail_mem_space_teams); } - // Add to list of mems. - mem->next = devr->memHead; - devr->memHead = mem; - *outMem = mem; free(globalSegmentInfo); return ret; fail_mem_space_teams: + { + struct ncclDevrMemory** ptr = &devr->memHead; + while (*ptr != nullptr && *ptr != mem) ptr = &(*ptr)->next; + if (*ptr == mem) *ptr = mem->next; + } + if (mem->ginSegmentInfos != nullptr) { + for (int segment = 0; segment < mem->numGinSegments; segment++) { + ncclGinDeregister(comm, mem->ginSegmentInfos[segment].ginHostWins); + } + } + if (mem->rmaHostWins[0] != nullptr) { + (void)ncclRmaProxyDeregister(comm, mem->rmaHostWins); + } for (struct ncclDevrTeam* t = devr->teamHead; t != nullptr; t = t->next) { symUnbindTeamMemory(comm, t, mem); if (ucBound && t->ucLeId != NCCL_LE_ID_INVALID) { @@ -888,6 +904,7 @@ static ncclResult_t symMemoryObtain(struct ncclComm* comm, CUmemGenericAllocatio ncclSpaceFree(&devr->bigSpace, bigOffset, mem->lsaMaxSize); fail_mem: if (mem != nullptr) { + free(mem->ginSegmentInfos); free(mem->memHandles); free(mem->segmentSizes); free(mem->lsaNumSegments); From 2dc1f41ff2925b1b8deec5883b1a844e67a03a03 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Fri, 4 Sep 2026 15:07:29 -0400 Subject: [PATCH 02/17] test(rccl): cover window register after ncclDevCommCreate (AICOMRCCL-2241) Alltoall_PureReference now registers before DevCommCreate, so it no longer catches the Anvil SDMA memHead lookup. Add a GIN-enabled case that does. Co-authored-by: Cursor --- .../rccl/test/transport/GinDeviceMPITests.cpp | 40 +++++++++++++++++++ .../gin_rocshmem_backends_mellanox.json | 1 + 2 files changed, 41 insertions(+) diff --git a/projects/rccl/test/transport/GinDeviceMPITests.cpp b/projects/rccl/test/transport/GinDeviceMPITests.cpp index f25988ffba2e..95e0636bd27d 100644 --- a/projects/rccl/test/transport/GinDeviceMPITests.cpp +++ b/projects/rccl/test/transport/GinDeviceMPITests.cpp @@ -2221,6 +2221,46 @@ TEST_F(GinMPIDeviceTests, BarrierSession_LsaOnly) { MPI_Barrier(MPI_COMM_WORLD); } +// AICOMRCCL-2241: registering a symmetric window after ncclDevCommCreate is +// legal. Anvil SDMA resolves the user VA via ncclDevrGetLsaSelfAddr, which +// only walks memHead, so ginEnabled must already see this mem on the list. +// BarrierSession_LsaOnly also registers after create, but it never activates +// GIN (plain NCCL_DEV_COMM_REQUIREMENTS_INITIALIZER), so it cannot catch this. +TEST_F(GinMPIDeviceTests, WindowRegister_AfterDevCommCreate) { + if (auto reason = ginProxyTestSkipReason(); !reason.empty()) + GTEST_SKIP() << reason; + + if (!validateTestPrerequisites(/*min_processes=*/2, /*max_processes=*/8)) + GTEST_SKIP() << "Requires 2-8 ranks"; + + ASSERT_EQ(ncclSuccess, createTestCommunicator()); + ncclComm_t comm = getActiveCommunicator(); + + ncclDevCommRequirements reqs = defaultGinReqs(); + reqs.ginSignalCount = 1; + ncclDevComm devComm{}; + ASSERT_MPI_EQ(ncclSuccess, ncclDevCommCreate(comm, &reqs, &devComm)); + auto devCommCleanup = makeScopeGuard([&]() { + (void)ncclDevCommDestroy(comm, &devComm); + }); + + constexpr size_t kBytes = 64; + void* dBuf = nullptr; + ASSERT_MPI_EQ(ncclSuccess, ncclMemAlloc(&dBuf, kBytes)); + auto memCleanup = makeScopeGuard([&]() { + if (dBuf) (void)ncclMemFree(dBuf); + }); + + ncclWindow_t win = nullptr; + ASSERT_MPI_EQ(ncclSuccess, + ncclCommWindowRegister(comm, dBuf, kBytes, &win, NCCL_WIN_COLL_SYMMETRIC)); + auto winCleanup = makeScopeGuard([&]() { + if (win) (void)ncclCommWindowDeregister(comm, win); + }); + + MPI_Barrier(MPI_COMM_WORLD); +} + // Collective over the world team: composes the inner LSA barrier with the // outer rail-GIN barrier (ncclTeamTagWorld ctor). Mirrors Barrier_TwoRanks' // "completion == success" philosophy -- a broken inner LSA epoch or outer diff --git a/projects/rccl/tools/scripts/test_runner/configs/gin_rocshmem_backends_mellanox.json b/projects/rccl/tools/scripts/test_runner/configs/gin_rocshmem_backends_mellanox.json index 6351893d2bf6..1b9f0b2653c8 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/gin_rocshmem_backends_mellanox.json +++ b/projects/rccl/tools/scripts/test_runner/configs/gin_rocshmem_backends_mellanox.json @@ -129,6 +129,7 @@ "name": "BarrierSession_LsaOnly", "description": "ncclBarrierSession LSA-only (2 ranks/node, non-trivial LSA team)", "test_filter": "GinMPIDeviceTests.BarrierSession_LsaOnly", "num_ranks": 4, "num_gpus": 2 }, + { "name": "WindowRegister_AfterDevCommCreate", "description": "Symmetric window register after ncclDevCommCreate (AICOMRCCL-2241)", "test_filter": "GinMPIDeviceTests.WindowRegister_AfterDevCommCreate" }, { "name": "BarrierSession_Hybrid", "description": "ncclBarrierSession hybrid inner-LSA + outer rail-GIN (2 ranks/node)", "test_filter": "GinMPIDeviceTests.BarrierSession_Hybrid", "num_ranks": 4, "num_gpus": 2 From 1fd6c121f7e48ba1d8c02ca80e9bab38601a840e Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Thu, 10 Sep 2026 22:26:43 -0400 Subject: [PATCH 03/17] test(rccl): run WindowRegister_AfterDevCommCreate on MI300/MI355 CI Name the AICOMRCCL-2241 case in the cluster test-runner configs so proxy device-MPI jobs report it explicitly, not only via the catch-all filter. Co-authored-by: Cursor --- .../scripts/test_runner/configs/mi300x_mellanox_ib.json | 5 +++++ .../tools/scripts/test_runner/configs/mi355x_thor2_roce.json | 5 +++++ 2 files changed, 10 insertions(+) diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json b/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json index 89c162fda247..4622fcaf781b 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json @@ -3958,6 +3958,11 @@ "test_filter": "GinMPIDeviceTests.*:-GinMPIDeviceTests.VASignal_*:-GinMPIDeviceTests.SignalAdd_PutValue:-GinMPIDeviceTests.SignalInc_MixingRule:-GinMPIDeviceTests.Signal_WaitRead_Low32Bits:-GinMPIDeviceTests.SignalShadow_GetPtr:-GinMPIDeviceTests.MultiContext_NonPowerOf2:-GinMPIDeviceTests.RailConnection_Create:-GinMPIDeviceTests.*_MultiContext", "timeout": 600 }, + { + "name": "GIN_WindowRegister_AfterDevCommCreate", + "description": "Symmetric window register after ncclDevCommCreate (AICOMRCCL-2241)", + "test_filter": "GinMPIDeviceTests.WindowRegister_AfterDevCommCreate" + }, { "name": "GIN_DeviceMPI_MultiContext", "description": "GIN device tests driven across multiple GIN contexts (NCCL_GIN_NCONTEXTS=4): put/putValue-inline/counter+signal/VA-signal, one block per context", diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json b/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json index ea39519ed276..7b787043ff69 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json @@ -4379,6 +4379,11 @@ "test_filter": "GinMPIDeviceTests.*", "timeout": 600 }, + { + "name": "GIN_WindowRegister_AfterDevCommCreate", + "description": "Symmetric window register after ncclDevCommCreate (AICOMRCCL-2241)", + "test_filter": "GinMPIDeviceTests.WindowRegister_AfterDevCommCreate" + }, { "name": "GIN_MultiContext_NonPowerOf2", "description": "Multi-context GIN with a non-power-of-2 context count (NCCL_GIN_NCONTEXTS=3)", From c23da96f7bc281e932469296c44b370f82fe66ca Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Fri, 11 Sep 2026 16:05:17 -0400 Subject: [PATCH 04/17] test(rccl): drop WindowRegister CI entries that cannot catch AICOMRCCL-2241 Those gin_device_mpi jobs pin NCCL_GIN_TYPE=2 on two nodes, so they never reach Anvil SDMA memHead lookup. Coverage stays on gin_devmpi_sdma. Co-authored-by: Cursor --- .../scripts/test_runner/configs/mi300x_mellanox_ib.json | 5 ----- .../tools/scripts/test_runner/configs/mi355x_thor2_roce.json | 5 ----- 2 files changed, 10 deletions(-) diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json b/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json index 4622fcaf781b..89c162fda247 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi300x_mellanox_ib.json @@ -3958,11 +3958,6 @@ "test_filter": "GinMPIDeviceTests.*:-GinMPIDeviceTests.VASignal_*:-GinMPIDeviceTests.SignalAdd_PutValue:-GinMPIDeviceTests.SignalInc_MixingRule:-GinMPIDeviceTests.Signal_WaitRead_Low32Bits:-GinMPIDeviceTests.SignalShadow_GetPtr:-GinMPIDeviceTests.MultiContext_NonPowerOf2:-GinMPIDeviceTests.RailConnection_Create:-GinMPIDeviceTests.*_MultiContext", "timeout": 600 }, - { - "name": "GIN_WindowRegister_AfterDevCommCreate", - "description": "Symmetric window register after ncclDevCommCreate (AICOMRCCL-2241)", - "test_filter": "GinMPIDeviceTests.WindowRegister_AfterDevCommCreate" - }, { "name": "GIN_DeviceMPI_MultiContext", "description": "GIN device tests driven across multiple GIN contexts (NCCL_GIN_NCONTEXTS=4): put/putValue-inline/counter+signal/VA-signal, one block per context", diff --git a/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json b/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json index 7b787043ff69..ea39519ed276 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json +++ b/projects/rccl/tools/scripts/test_runner/configs/mi355x_thor2_roce.json @@ -4379,11 +4379,6 @@ "test_filter": "GinMPIDeviceTests.*", "timeout": 600 }, - { - "name": "GIN_WindowRegister_AfterDevCommCreate", - "description": "Symmetric window register after ncclDevCommCreate (AICOMRCCL-2241)", - "test_filter": "GinMPIDeviceTests.WindowRegister_AfterDevCommCreate" - }, { "name": "GIN_MultiContext_NonPowerOf2", "description": "Multi-context GIN with a non-power-of-2 context count (NCCL_GIN_NCONTEXTS=3)", From 91f783d0fcc6ff48020e8e46b5fc1685a77bae3c Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Sat, 12 Sep 2026 03:35:12 -0400 Subject: [PATCH 05/17] fix(rccl): address Argus review on symMemoryObtain unwind (AICOMRCCL-2241) Route symMemoryRegisterGin failures through NCCLCHECKGOTO with a fail_simple path that deregisters partial MRs, factor GIN/RMA teardown into symMemoryUnregister, and add a DevRuntime unit test that asserts memHead is unlinked when ncclGinRegister fails after the pre-register link. Co-authored-by: Cursor --- projects/rccl/src/dev_runtime.cc | 115 ++++++++++----------- projects/rccl/test/DevRuntimeTestsStubs.cc | 4 + projects/rccl/test/gin_teardown_test.cpp | 18 ++++ 3 files changed, 77 insertions(+), 60 deletions(-) diff --git a/projects/rccl/src/dev_runtime.cc b/projects/rccl/src/dev_runtime.cc index cc50a9c36ab7..37ef893cb15b 100644 --- a/projects/rccl/src/dev_runtime.cc +++ b/projects/rccl/src/dev_runtime.cc @@ -713,15 +713,14 @@ static ncclResult_t symMemoryRegisterGin(struct ncclComm* comm, struct ncclDevrM mem->ginHostWins[i] = mem->ginSegmentInfos[0].ginHostWins[i]; } } -exit: return ret; fail: for (int i = 0; i < numSegmentsRegistered; i++) { - ncclGinDeregister(comm, mem->ginSegmentInfos[i].ginHostWins); + (void)ncclGinDeregister(comm, mem->ginSegmentInfos[i].ginHostWins); } free(mem->ginSegmentInfos); mem->ginSegmentInfos = nullptr; - goto exit; + return ret; } static ncclResult_t symMemoryRegisterRma(struct ncclComm* comm, struct ncclDevrMemory* mem) { @@ -732,6 +731,18 @@ static ncclResult_t symMemoryRegisterRma(struct ncclComm* comm, struct ncclDevrM return ncclSuccess; } +static void symMemoryUnregister(struct ncclComm* comm, struct ncclDevrMemory* mem) { + struct ncclDevrState* devr = &comm->devrState; + if (devr->ginEnabled && mem->ginSegmentInfos != nullptr) { + for (int segment = 0; segment < mem->numGinSegments; segment++) { + ncclGinDeregister(comm, mem->ginSegmentInfos[segment].ginHostWins); + } + } + if (devr->rmaProxyEnabled && mem->maxGlobalNumSegments == 1) { + (void)ncclRmaProxyDeregister(comm, mem->rmaHostWins); + } +} + // On success we take caller's reference on memHandle. // Due to multicast binds for each pre-exiting team, this function requires // caller do a world barrier before returning to user. @@ -878,19 +889,12 @@ static ncclResult_t symMemoryObtain(struct ncclComm* comm, CUmemGenericAllocatio return ret; fail_mem_space_teams: + symMemoryUnregister(comm, mem); { struct ncclDevrMemory** ptr = &devr->memHead; while (*ptr != nullptr && *ptr != mem) ptr = &(*ptr)->next; if (*ptr == mem) *ptr = mem->next; } - if (mem->ginSegmentInfos != nullptr) { - for (int segment = 0; segment < mem->numGinSegments; segment++) { - ncclGinDeregister(comm, mem->ginSegmentInfos[segment].ginHostWins); - } - } - if (mem->rmaHostWins[0] != nullptr) { - (void)ncclRmaProxyDeregister(comm, mem->rmaHostWins); - } for (struct ncclDevrTeam* t = devr->teamHead; t != nullptr; t = t->next) { symUnbindTeamMemory(comm, t, mem); if (ucBound && t->ucLeId != NCCL_LE_ID_INVALID) { @@ -916,59 +920,50 @@ static ncclResult_t symMemoryObtain(struct ncclComm* comm, CUmemGenericAllocatio } static void symMemoryDestroy(struct ncclComm* comm, struct ncclDevrMemory* mem) { - if (mem != nullptr) { - struct ncclDevrState* devr = &comm->devrState; - // Idempotent: a window pair or the finalize drain may already have destroyed - // this mem. Establish membership before touching it, not after. The unlink - // below walks the list unguarded, so a second call on a freed mem would - // otherwise repeat every deregister, unmap, cuMemRelease and free in this - // function and then run off the end of the list. - { - struct ncclDevrMemory* cursor = devr->memHead; - while (cursor != nullptr && cursor != mem) cursor = cursor->next; - if (cursor != mem) return; - } - if (devr->ginEnabled && mem->ginSegmentInfos != nullptr) { - for (int segment = 0; segment < mem->numGinSegments; segment++) { - ncclGinDeregister(comm, mem->ginSegmentInfos[segment].ginHostWins); - } - } - if (devr->rmaProxyEnabled && mem->maxGlobalNumSegments == 1) { - ncclRmaProxyDeregister(comm, mem->rmaHostWins); - } - ncclCftLeId leUcSelf = devr->le.baseId == NCCL_LE_ID_INVALID ? NCCL_LE_ID_INVALID : devr->le.baseId + devr->cftSelf; - symUnbindTeamLe(comm, mem, leUcSelf); - for (struct ncclDevrTeam* t = devr->teamHead; t != nullptr; t = t->next) { - symUnbindTeamMemory(comm, t, mem); - symUnbindTeamLe(comm, mem, t->mcLeId); - } - for (int r = 0; r < devr->lsaSize; r++) { - uintptr_t base = reinterpret_cast(devr->lsaFlatBase); - uintptr_t addr = base + r * devr->bigSize + mem->bigOffset; - for (int idx = 0; idx < mem->lsaNumSegments[r]; idx++) { - CUdeviceptr tmpBase; - size_t tmpBaseSize; - CUCHECKIGNORE(cuMemGetAddressRange(&tmpBase, &tmpBaseSize, reinterpret_cast(addr))); - CUCHECKIGNORE(cuMemUnmap(reinterpret_cast(addr), tmpBaseSize)); - addr = addr + tmpBaseSize; - } - } + if (mem == nullptr) { + return; + } + struct ncclDevrState* devr = &comm->devrState; + struct ncclDevrMemory** memLink = &devr->memHead; + while (*memLink != nullptr && *memLink != mem) { + memLink = &(*memLink)->next; + } + // Idempotent: a window pair or finalize drain may already have destroyed this mem. + if (*memLink != mem) { + return; + } - ncclSpaceFree(&devr->bigSpace, mem->bigOffset, mem->lsaMaxSize); - for (int segment = 0; segment < mem->numSegments; segment++) { - CUCHECKIGNORE(cuMemRelease(mem->memHandles[segment])); + symMemoryUnregister(comm, mem); + ncclCftLeId leUcSelf = devr->le.baseId == NCCL_LE_ID_INVALID ? NCCL_LE_ID_INVALID : devr->le.baseId + devr->cftSelf; + symUnbindTeamLe(comm, mem, leUcSelf); + for (struct ncclDevrTeam* t = devr->teamHead; t != nullptr; t = t->next) { + symUnbindTeamMemory(comm, t, mem); + symUnbindTeamLe(comm, mem, t->mcLeId); + } + for (int r = 0; r < devr->lsaSize; r++) { + uintptr_t base = reinterpret_cast(devr->lsaFlatBase); + uintptr_t addr = base + r * devr->bigSize + mem->bigOffset; + for (int idx = 0; idx < mem->lsaNumSegments[r]; idx++) { + CUdeviceptr tmpBase; + size_t tmpBaseSize; + CUCHECKIGNORE(cuMemGetAddressRange(&tmpBase, &tmpBaseSize, reinterpret_cast(addr))); + CUCHECKIGNORE(cuMemUnmap(reinterpret_cast(addr), tmpBaseSize)); + addr = addr + tmpBaseSize; } + } - struct ncclDevrMemory** ptr = &devr->memHead; - while (*ptr != mem) ptr = &(*ptr)->next; - *ptr = mem->next; // Remove from list. - - free(mem->ginSegmentInfos); - free(mem->lsaNumSegments); - free(mem->segmentSizes); - free(mem->memHandles); - free(mem); + ncclSpaceFree(&devr->bigSpace, mem->bigOffset, mem->lsaMaxSize); + for (int segment = 0; segment < mem->numSegments; segment++) { + CUCHECKIGNORE(cuMemRelease(mem->memHandles[segment])); } + + *memLink = mem->next; // Remove from list. + + free(mem->ginSegmentInfos); + free(mem->lsaNumSegments); + free(mem->segmentSizes); + free(mem->memHandles); + free(mem); } static ncclResult_t symWindowTableInitOnce(struct ncclComm* comm, cudaStream_t stream) { diff --git a/projects/rccl/test/DevRuntimeTestsStubs.cc b/projects/rccl/test/DevRuntimeTestsStubs.cc index bf6777ec81e7..1e8558d10a69 100644 --- a/projects/rccl/test/DevRuntimeTestsStubs.cc +++ b/projects/rccl/test/DevRuntimeTestsStubs.cc @@ -185,6 +185,9 @@ void* ncclMemoryStack::allocateSpilled(struct ncclMemoryStack*, size_t size, siz // --------------------------------------------------------------------------- // GIN host. // --------------------------------------------------------------------------- +static int devRuntimeTestGinRegisterFail = 0; + +extern "C" void DevRuntimeTests_SetGinRegisterFail(int fail) { devRuntimeTestGinRegisterFail = fail; } ncclResult_t ncclGetGinType(struct ncclComm*, ncclGinType_t* ginType) { if (ginType) *ginType = NCCL_GIN_TYPE_NONE; return ncclSuccess; @@ -201,6 +204,7 @@ ncclResult_t ncclGinDevCommSetup(struct ncclComm*, struct ncclDevCommRequirement ncclResult_t ncclGinDevCommFree(struct ncclComm*, struct ncclDevComm const*) { return ncclSuccess; } ncclResult_t ncclGinRegister(struct ncclComm*, void*, size_t, void*[NCCL_GIN_MAX_CONNECTIONS], ncclGinWindow_t[NCCL_GIN_MAX_CONNECTIONS], int, bool, int) { + if (devRuntimeTestGinRegisterFail) return ncclInternalError; return ncclSuccess; } ncclResult_t ncclGinDeregister(struct ncclComm*, void*[NCCL_GIN_MAX_CONNECTIONS]) { return ncclSuccess; } diff --git a/projects/rccl/test/gin_teardown_test.cpp b/projects/rccl/test/gin_teardown_test.cpp index 381969e99c0f..a0a118b294ad 100644 --- a/projects/rccl/test/gin_teardown_test.cpp +++ b/projects/rccl/test/gin_teardown_test.cpp @@ -19,6 +19,7 @@ #include "common/ProcessIsolatedTestRunner.hpp" extern int rcclTestHipMemAddressFreeCount; +extern "C" void DevRuntimeTests_SetGinRegisterFail(int fail); // Build the smallest ncclComm/ncclDevrState that symMemoryObtain will accept: // a single-rank, single-LSA-team comm with GIN and RMA proxy disabled. @@ -111,6 +112,23 @@ TEST_F(SymMemoryObtainTest, DestroyIsIdempotentWhileAnotherMemoryIsLinked) { EXPECT_EQ(comm->devrState.memHead, nullptr); } +// symMemoryObtain links mem onto memHead before GIN registration. If registration +// fails after that point, fail_mem_space_teams must unlink it again. +TEST_F(SymMemoryObtainTest, ObtainFailureUnlinksMemHeadWhenGinRegisterFails) { + comm->devrState.ginEnabled = true; + DevRuntimeTests_SetGinRegisterFail(1); + + hipMemGenericAllocationHandle_t memHandle = reinterpret_cast(0x1); + void* userAddr = reinterpret_cast(0x100000); + struct ncclDevrMemory* mem = nullptr; + + EXPECT_NE(symMemoryObtain(comm, &memHandle, /*numSegments=*/1, userAddr, /*size=*/4096, /*winFlags=*/0, &mem), + ncclSuccess); + EXPECT_EQ(comm->devrState.memHead, nullptr); + + DevRuntimeTests_SetGinRegisterFail(0); +} + // --------------------------------------------------------------------------- // AICOMRCCL-835 finalize-drain coverage. // From 69d47b3192b9c13fb16e88f0f37d2b384f1874ab Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Sat, 12 Sep 2026 03:43:00 -0400 Subject: [PATCH 06/17] fix(rccl): skip null RMA handles on proxy deregister and note AICOMRCCL-2241 ncclRmaProxyRegister can return after filling only a prefix of rmaHostWins; deregMrSym implementations deref the handle, so skip null slots in ncclRmaProxyDeregister. Record the Anvil SDMA window-register fix in CHANGELOG. Co-authored-by: Cursor --- projects/rccl/CHANGELOG.md | 1 + projects/rccl/src/rma/rma_proxy.cc | 4 ++++ 2 files changed, 5 insertions(+) diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md index ad9a3f5164dd..4b67b60f0e28 100644 --- a/projects/rccl/CHANGELOG.md +++ b/projects/rccl/CHANGELOG.md @@ -63,6 +63,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https: * Fixed gfx1250 LL and LL128 comm-FIFO hangs on sibling DPX partitions. The FIFO store now uses system-scope b128 (`RCCL_LL_FIFO_SYS_SCOPE`) alongside the existing system-scope load, preventing hangs at slot reuse starting from the 9th collective operation. * Fixed DDA fabric AllToAll validation race by staging send data into scratch with a host-launched `cudaMemcpyAsync` before the peer exchange kernel. * Fixed DDA fabric barrier publication race causing sporadic validation errors by using release-acquire semantics on the prologue barrier. +* Fixed `ncclCommWindowRegister` after `ncclDevCommCreate` aborting on the GIN Anvil SDMA backend (`NCCL_GIN_TYPE=6`) with `could not resolve LSA flat addr`. Symmetric-window memory is now linked onto the device-runtime list before GIN/RMA registration so plugins that resolve the user VA through that list can see the in-flight window. ### Known issues * On gfx90a (MI210/MI250/MI250X) with ROCm 7.13 or later, per-launch scratch-memory reclaim in the runtime degrades RCCL performance. Set `HSA_NO_SCRATCH_RECLAIM=1` to restore performance. diff --git a/projects/rccl/src/rma/rma_proxy.cc b/projects/rccl/src/rma/rma_proxy.cc index 8d418aa6d968..f1a389a61af7 100644 --- a/projects/rccl/src/rma/rma_proxy.cc +++ b/projects/rccl/src/rma/rma_proxy.cc @@ -410,6 +410,10 @@ ncclResult_t ncclRmaProxyRegister(struct ncclComm* comm, void* address, size_t s ncclResult_t ncclRmaProxyDeregister(struct ncclComm* comm, void* rmaHostWins[NCCL_RMA_MAX_CONNECTIONS]) { struct ncclRmaProxyState* rmaProxyState = &comm->rmaState.rmaProxyState; for (int n = 0; n < rmaProxyState->rmaCommCount; n++) { + // ncclRmaProxyRegister returns at the first failing index, so later slots + // (and, on a partial success, any unused connections) stay null. deregMrSym + // implementations deref the handle without a null check. + if (rmaHostWins[n] == nullptr) continue; NCCLCHECK(rmaProxyState->ncclRma->deregMrSym(rmaProxyState->rmaComms[n], rmaHostWins[n])); } return ncclSuccess; From f4383395ca8330462d87427d39b830352a9c73fd Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Sat, 12 Sep 2026 03:48:24 -0400 Subject: [PATCH 07/17] test(rccl): add SDMA-only GinMPIDeviceTests runner config (AICOMRCCL-2241) gin_rocshmem_backends_mellanox.json also runs GDA device-MPI and alltoall_perf. Add gin_devmpi_sdma.json so this PR can trigger only the Anvil SDMA suite. Co-authored-by: Cursor --- .../test_runner/configs/gin_devmpi_sdma.json | 152 ++++++++++++++++++ 1 file changed, 152 insertions(+) create mode 100644 projects/rccl/tools/scripts/test_runner/configs/gin_devmpi_sdma.json diff --git a/projects/rccl/tools/scripts/test_runner/configs/gin_devmpi_sdma.json b/projects/rccl/tools/scripts/test_runner/configs/gin_devmpi_sdma.json new file mode 100644 index 000000000000..15f30078571d --- /dev/null +++ b/projects/rccl/tools/scripts/test_runner/configs/gin_devmpi_sdma.json @@ -0,0 +1,152 @@ +{ + "system_configurations": { + "name": "RCCL-GIN-Device-MPI-SDMA-MI300X-Mellanox-IB", + "description": "GinMPIDeviceTests on the Anvil SDMA backend only (NCCL_GIN_TYPE=6, single node). No GDA and no alltoall_perf." + }, + "paths": { + "workdir": "${WORKDIR:-$PWD}", + "rocm_path": "${ROCM_PATH:-/opt/rocm}", + "mpi_path": "${MPI_PATH:-/opt/shared/ompi}" + }, + "env_variables": { + "NCCL_SOCKET_IFNAME": "eth0,eth1", + "HSA_NO_SCRATCH_RECLAIM": "1", + "NCCL_DEBUG": "INFO" + }, + "build_configuration": { + "install_flags": [ + "-l", + "--debug", + "-t", + "--enable-mpi-tests", + "--enable-code-coverage", + "--rocshmem-gin", + "--prefix", + "${WORKDIR}/install" + ], + "parallel_jobs": 64 + }, + "rccl_tests_build_configuration": { + "enabled": false + }, + "test_configurations": { + "default": { + "env_variables": {}, + "rerun_env_variables": { + "NCCL_DEBUG": "INFO" + } + }, + "gin_devmpi_common": { + "extends": "default", + "is_gtest": true, + "binary": "rccl-UnitTestsMPI", + "num_ranks": 2, + "num_nodes": 1, + "num_gpus": 2, + "timeout": 180, + "env_variables": { + "NCCL_GIN_ENABLE": "1", + "NCCL_CUMEM_ENABLE": "1", + "NCCL_ENV_PLUGIN": "none", + "NCCL_DMABUF_ENABLE": "1" + }, + "tests": [ + { "name": "Put_BasicAndOffsets", "description": "GIN device put with base + offset variants", "test_filter": "GinMPIDeviceTests.Put_BasicAndOffsets" }, + { "name": "Put_CrossNode", "description": "GIN device put crossing the IB fabric", "test_filter": "GinMPIDeviceTests.Put_CrossNode" }, + { "name": "PutValue_Inline", "description": "GIN device inline put-value", "test_filter": "GinMPIDeviceTests.PutValue_Inline" }, + { "name": "Signal_NoPayload", "description": "GIN device signal, no payload", "test_filter": "GinMPIDeviceTests.Signal_NoPayload" }, + { "name": "Signal_HighIdNoOverflow", "description": "GIN device signal with high id, no overflow", "test_filter": "GinMPIDeviceTests.Signal_HighIdNoOverflow" }, + { "name": "WaitCounter_Local", "description": "GIN device wait-counter, local", "test_filter": "GinMPIDeviceTests.WaitCounter_Local" }, + { "name": "WaitCounterAndSignal", "description": "GIN device wait-counter combined with signal", "test_filter": "GinMPIDeviceTests.WaitCounterAndSignal" }, + { "name": "Barrier_TwoRanks", "description": "GIN device two-rank barrier", "test_filter": "GinMPIDeviceTests.Barrier_TwoRanks" }, + { + "name": "BarrierSession_LsaOnly", "description": "ncclBarrierSession LSA-only (2 ranks/node, non-trivial LSA team)", "test_filter": "GinMPIDeviceTests.BarrierSession_LsaOnly", + "num_ranks": 4, "num_gpus": 4 + }, + { "name": "WindowRegister_AfterDevCommCreate", "description": "Symmetric window register after ncclDevCommCreate (AICOMRCCL-2241)", "test_filter": "GinMPIDeviceTests.WindowRegister_AfterDevCommCreate" }, + { + "name": "BarrierSession_Hybrid", "description": "ncclBarrierSession hybrid inner-LSA + outer rail-GIN (2 ranks/node)", "test_filter": "GinMPIDeviceTests.BarrierSession_Hybrid", + "num_ranks": 4, "num_gpus": 4 + }, + { "name": "Barrier_WorldMultiIndex", "description": "World-team GIN barrier, 4 concurrent barrier indices", "test_filter": "GinMPIDeviceTests.Barrier_WorldMultiIndex" }, + { + "name": "BarrierPools_NoCrossTalk", "description": "Generic and specialized GIN barrier sessions keep separate signal pools (2 ranks/node, self-skips on a single node)", "test_filter": "GinMPIDeviceTests.BarrierPools_NoCrossTalk", + "num_ranks": 4, "num_gpus": 4 + }, + { "name": "SignalAdd_AndShadow", "description": "GIN device signal-add and shadow counter", "test_filter": "GinMPIDeviceTests.SignalAdd_AndShadow" }, + { "name": "SymPtr_PutAndPutValue", "description": "GIN symmetric-pointer put and put-value", "test_filter": "GinMPIDeviceTests.SymPtr_PutAndPutValue" }, + { "name": "Alltoall_PureReference", "description": "GIN pure-reference alltoall kernel", "test_filter": "GinMPIDeviceTests.Alltoall_PureReference" }, + { "name": "Alltoall_WorldGinBarrier", "description": "GIN alltoall on the world-team GIN barrier pool", "test_filter": "GinMPIDeviceTests.Alltoall_WorldGinBarrierReference" }, + { + "name": "AlltoallHybrid_Reference", "description": "GIN hybrid (LSA + GIN) alltoall reference (2 ranks/node)", "test_filter": "GinMPIDeviceTests.AlltoallHybrid_Reference", + "num_ranks": 4, "num_gpus": 4 + }, + { + "name": "Alltoall_SdmaInternal", "description": "GIN-SDMA alltoall via ncclAlltoAll (SDMA only, 8 ranks)", "test_filter": "GinMPIDeviceTests.Alltoall_SdmaInternal", + "num_ranks": 8, "num_nodes": 1, "num_gpus": 8 + }, + { "name": "MultiContext_AllFourRoute","description": "GIN multi-context, all four routes", "test_filter": "GinMPIDeviceTests.MultiContext_AllFourRoute" }, + { + "name": "MultiContext_NonPowerOf2", "description": "GIN multi-context with non-power-of-2 context count (NCONNECTIONS=3, NCONTEXTS=3)", "test_filter": "GinMPIDeviceTests.MultiContext_NonPowerOf2", + "env_variables": { "NCCL_GIN_NCONNECTIONS": "3", "NCCL_GIN_NCONTEXTS": "3" } + }, + { "name": "LargeBuffer_Sweep", "description": "GIN large-buffer size sweep", "test_filter": "GinMPIDeviceTests.LargeBuffer_Sweep", "timeout": 300 }, + { + "name": "Disable_Error", "description": "Negative: devCommCreate must reject GIN when NCCL_GIN_ENABLE=0", "test_filter": "GinMPIDeviceTests.Disable_Error", + "env_variables": { "NCCL_GIN_ENABLE": "0" } + }, + { + "name": "Invalid_SignalPool", "description": "Negative: oversized signal pool rejected at devCommCreate", "test_filter": "GinMPIDeviceTests.Invalid_SignalPool", + "env_variables": { "NCCL_GIN_SIGNAL_POOL_SIZE": "0x40000000" } + }, + { + "name": "Invalid_CounterPool", "description": "Negative: oversized counter pool rejected at devCommCreate", "test_filter": "GinMPIDeviceTests.Invalid_CounterPool", + "env_variables": { "NCCL_GIN_COUNTER_POOL_SIZE": "0x40000000" } + }, + { "name": "Teardown_NoLeaks", "description": "GIN teardown leaves no leaks", "test_filter": "GinMPIDeviceTests.Teardown_NoLeaks" }, + { "name": "Init_Destroy_Stress", "description": "GIN init/destroy stress loop", "test_filter": "GinMPIDeviceTests.Init_Destroy_Stress", "timeout": 300 }, + { "name": "Alltoall_CrossNode", "description": "GIN cross-node alltoall (self-skips unless >=16 ranks)", "test_filter": "GinMPIDeviceTests.Alltoall_CrossNode" }, + { "name": "Properties_NLsaTeams", "description": "GIN properties: number of LSA teams", "test_filter": "GinMPIDeviceTests.Properties_NLsaTeams" }, + { "name": "MultiContext_Exclusive", "description": "GIN multi-context exclusive contexts", "test_filter": "GinMPIDeviceTests.MultiContext_Exclusive" }, + { "name": "VASignal_Put", "description": "GIN VA-signal put", "test_filter": "GinMPIDeviceTests.VASignal_Put" }, + { + "name": "RailConnection_Create", "description": "GIN rail connection create; rail-only mode (NCCL_CROSS_NIC=0)", "test_filter": "GinMPIDeviceTests.RailConnection_Create", + "env_variables": { "NCCL_CROSS_NIC": "0" } + }, + { + "name": "ReduceScatter_Symmetric", "description": "Symmetric-memory ReduceScatter sum (2 ranks/node)", "test_filter": "GinMPIDeviceTests.ReduceScatter_Symmetric", + "num_ranks": 4, "num_gpus": 4 + }, + { + "name": "ReduceScatter_Symmetric_Avg", "description": "Symmetric-memory ReduceScatter avg (2 ranks/node)", "test_filter": "GinMPIDeviceTests.ReduceScatter_Symmetric_Avg", + "num_ranks": 4, "num_gpus": 4 + } + ] + }, + "gin_devmpi_sdma": { + "extends": "gin_devmpi_common", + "description": "SDMA is a single-node backend: every GinMPIDeviceTests case runs on a single node (num_nodes=1). Cross-node cases (Put_CrossNode/Alltoall_CrossNode/ReduceScatter_Symmetric*) self-skip.", + "num_nodes": 1, + "num_gpus": 2, + "env_variables": { + "NCCL_GIN_TYPE": "6" + }, + "tests": [ + { "name": "BarrierSession_LsaOnly", "num_gpus": 4 }, + { "name": "BarrierSession_Hybrid", "num_gpus": 4 }, + { "name": "BarrierPools_NoCrossTalk", "num_gpus": 4 }, + { "name": "AlltoallHybrid_Reference", "num_gpus": 4 }, + { "name": "ReduceScatter_Symmetric", "num_gpus": 4 }, + { "name": "ReduceScatter_Symmetric_Avg", "num_gpus": 4 } + ] + } + }, + "test_suites": [ + { + "name": "GIN Device MPI Tests - SDMA (single node)", + "description": "Every GinMPIDeviceTests case, one mpirun per test, on the SDMA backend (single node)", + "config": "gin_devmpi_sdma", + "enabled": true + } + ] +} From 7ecc873b8d67ff2f194040b41df8f4e3b78f4893 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Sat, 12 Sep 2026 13:21:18 -0400 Subject: [PATCH 08/17] fix(rccl): skip null GIN handles on deregister and pin memHead ordering (AICOMRCCL-2241) fail_simple can hand a partly filled ginHostWins to ncclGinDeregister; skip null slots the way ncclRmaProxyDeregister already does. Record memHead at ncclGinRegister so the host test fails if the pre-register link moves. Keep gin_devmpi_sdma.json as a five-test Anvil SDMA subset instead of the full suite that self-skips. Co-authored-by: Cursor --- projects/rccl/src/gin/gin_host.cc | 4 + projects/rccl/test/DevRuntimeTestsStubs.cc | 10 +- projects/rccl/test/gin_teardown_test.cpp | 2 + .../test_runner/configs/gin_devmpi_sdma.json | 108 +++++------------- 4 files changed, 40 insertions(+), 84 deletions(-) diff --git a/projects/rccl/src/gin/gin_host.cc b/projects/rccl/src/gin/gin_host.cc index ac0ef27cfc76..a1fb1990fbfe 100644 --- a/projects/rccl/src/gin/gin_host.cc +++ b/projects/rccl/src/gin/gin_host.cc @@ -547,6 +547,10 @@ ncclResult_t ncclGinDeregister(struct ncclComm* comm, struct ncclGinBackendState* backend = &ginState->backends[backendIdx]; for (int commIdx = 0; commIdx < backend->ginCommCount; commIdx++) { int slot = backendIdx * NCCL_GIN_MAX_CONNECTIONS + commIdx; + // ncclGinRegister returns at the first failing index, and the multi-segment + // DMABUF bail leaves every slot null. deregMrSym implementations deref the + // handle without a null check. + if (ginHostWins[slot] == nullptr) continue; NCCLCHECK(backend->ncclGin->deregMrSym(backend->ginComms[commIdx], ginHostWins[slot])); } } diff --git a/projects/rccl/test/DevRuntimeTestsStubs.cc b/projects/rccl/test/DevRuntimeTestsStubs.cc index 1e8558d10a69..9071a84640ab 100644 --- a/projects/rccl/test/DevRuntimeTestsStubs.cc +++ b/projects/rccl/test/DevRuntimeTestsStubs.cc @@ -186,8 +186,13 @@ void* ncclMemoryStack::allocateSpilled(struct ncclMemoryStack*, size_t size, siz // GIN host. // --------------------------------------------------------------------------- static int devRuntimeTestGinRegisterFail = 0; +static struct ncclDevrMemory* ginRegisterMemHeadAtCall = nullptr; -extern "C" void DevRuntimeTests_SetGinRegisterFail(int fail) { devRuntimeTestGinRegisterFail = fail; } +extern "C" void DevRuntimeTests_SetGinRegisterFail(int fail) { + devRuntimeTestGinRegisterFail = fail; + ginRegisterMemHeadAtCall = nullptr; +} +extern "C" struct ncclDevrMemory* DevRuntimeTests_GinRegisterMemHeadAtCall() { return ginRegisterMemHeadAtCall; } ncclResult_t ncclGetGinType(struct ncclComm*, ncclGinType_t* ginType) { if (ginType) *ginType = NCCL_GIN_TYPE_NONE; return ncclSuccess; @@ -202,8 +207,9 @@ ncclResult_t ncclGinDevCommSetup(struct ncclComm*, struct ncclDevCommRequirement return ncclSuccess; } ncclResult_t ncclGinDevCommFree(struct ncclComm*, struct ncclDevComm const*) { return ncclSuccess; } -ncclResult_t ncclGinRegister(struct ncclComm*, void*, size_t, void*[NCCL_GIN_MAX_CONNECTIONS], +ncclResult_t ncclGinRegister(struct ncclComm* comm, void*, size_t, void*[NCCL_GIN_MAX_CONNECTIONS], ncclGinWindow_t[NCCL_GIN_MAX_CONNECTIONS], int, bool, int) { + ginRegisterMemHeadAtCall = comm ? comm->devrState.memHead : nullptr; if (devRuntimeTestGinRegisterFail) return ncclInternalError; return ncclSuccess; } diff --git a/projects/rccl/test/gin_teardown_test.cpp b/projects/rccl/test/gin_teardown_test.cpp index a0a118b294ad..2555fc353205 100644 --- a/projects/rccl/test/gin_teardown_test.cpp +++ b/projects/rccl/test/gin_teardown_test.cpp @@ -20,6 +20,7 @@ extern int rcclTestHipMemAddressFreeCount; extern "C" void DevRuntimeTests_SetGinRegisterFail(int fail); +extern "C" struct ncclDevrMemory* DevRuntimeTests_GinRegisterMemHeadAtCall(); // Build the smallest ncclComm/ncclDevrState that symMemoryObtain will accept: // a single-rank, single-LSA-team comm with GIN and RMA proxy disabled. @@ -124,6 +125,7 @@ TEST_F(SymMemoryObtainTest, ObtainFailureUnlinksMemHeadWhenGinRegisterFails) { EXPECT_NE(symMemoryObtain(comm, &memHandle, /*numSegments=*/1, userAddr, /*size=*/4096, /*winFlags=*/0, &mem), ncclSuccess); + EXPECT_NE(DevRuntimeTests_GinRegisterMemHeadAtCall(), nullptr); EXPECT_EQ(comm->devrState.memHead, nullptr); DevRuntimeTests_SetGinRegisterFail(0); diff --git a/projects/rccl/tools/scripts/test_runner/configs/gin_devmpi_sdma.json b/projects/rccl/tools/scripts/test_runner/configs/gin_devmpi_sdma.json index 15f30078571d..639ac99edab5 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/gin_devmpi_sdma.json +++ b/projects/rccl/tools/scripts/test_runner/configs/gin_devmpi_sdma.json @@ -1,7 +1,7 @@ { "system_configurations": { "name": "RCCL-GIN-Device-MPI-SDMA-MI300X-Mellanox-IB", - "description": "GinMPIDeviceTests on the Anvil SDMA backend only (NCCL_GIN_TYPE=6, single node). No GDA and no alltoall_perf." + "description": "Anvil SDMA (NCCL_GIN_TYPE=6) GinMPIDeviceTests that catch post-ncclDevCommCreate window register. Single node, no GDA, no alltoall_perf." }, "paths": { "workdir": "${WORKDIR:-$PWD}", @@ -36,8 +36,9 @@ "NCCL_DEBUG": "INFO" } }, - "gin_devmpi_common": { + "gin_devmpi_sdma": { "extends": "default", + "description": "AICOMRCCL-2241 Anvil SDMA subset: WindowRegister_AfterDevCommCreate plus the device-API cases that also register after ncclDevCommCreate. Omits cross-node, VA-signal, and Anvil ReduceScatter cases that self-skip at NCCL_GIN_TYPE=6 on one node.", "is_gtest": true, "binary": "rccl-UnitTestsMPI", "num_ranks": 2, @@ -48,103 +49,46 @@ "NCCL_GIN_ENABLE": "1", "NCCL_CUMEM_ENABLE": "1", "NCCL_ENV_PLUGIN": "none", - "NCCL_DMABUF_ENABLE": "1" + "NCCL_DMABUF_ENABLE": "1", + "NCCL_GIN_TYPE": "6" }, "tests": [ - { "name": "Put_BasicAndOffsets", "description": "GIN device put with base + offset variants", "test_filter": "GinMPIDeviceTests.Put_BasicAndOffsets" }, - { "name": "Put_CrossNode", "description": "GIN device put crossing the IB fabric", "test_filter": "GinMPIDeviceTests.Put_CrossNode" }, - { "name": "PutValue_Inline", "description": "GIN device inline put-value", "test_filter": "GinMPIDeviceTests.PutValue_Inline" }, - { "name": "Signal_NoPayload", "description": "GIN device signal, no payload", "test_filter": "GinMPIDeviceTests.Signal_NoPayload" }, - { "name": "Signal_HighIdNoOverflow", "description": "GIN device signal with high id, no overflow", "test_filter": "GinMPIDeviceTests.Signal_HighIdNoOverflow" }, - { "name": "WaitCounter_Local", "description": "GIN device wait-counter, local", "test_filter": "GinMPIDeviceTests.WaitCounter_Local" }, - { "name": "WaitCounterAndSignal", "description": "GIN device wait-counter combined with signal", "test_filter": "GinMPIDeviceTests.WaitCounterAndSignal" }, - { "name": "Barrier_TwoRanks", "description": "GIN device two-rank barrier", "test_filter": "GinMPIDeviceTests.Barrier_TwoRanks" }, - { - "name": "BarrierSession_LsaOnly", "description": "ncclBarrierSession LSA-only (2 ranks/node, non-trivial LSA team)", "test_filter": "GinMPIDeviceTests.BarrierSession_LsaOnly", - "num_ranks": 4, "num_gpus": 4 - }, - { "name": "WindowRegister_AfterDevCommCreate", "description": "Symmetric window register after ncclDevCommCreate (AICOMRCCL-2241)", "test_filter": "GinMPIDeviceTests.WindowRegister_AfterDevCommCreate" }, { - "name": "BarrierSession_Hybrid", "description": "ncclBarrierSession hybrid inner-LSA + outer rail-GIN (2 ranks/node)", "test_filter": "GinMPIDeviceTests.BarrierSession_Hybrid", - "num_ranks": 4, "num_gpus": 4 + "name": "BarrierSession_LsaOnly", + "description": "ncclBarrierSession LSA-only (2 ranks/node, non-trivial LSA team)", + "test_filter": "GinMPIDeviceTests.BarrierSession_LsaOnly", + "num_ranks": 4, + "num_gpus": 4 }, - { "name": "Barrier_WorldMultiIndex", "description": "World-team GIN barrier, 4 concurrent barrier indices", "test_filter": "GinMPIDeviceTests.Barrier_WorldMultiIndex" }, { - "name": "BarrierPools_NoCrossTalk", "description": "Generic and specialized GIN barrier sessions keep separate signal pools (2 ranks/node, self-skips on a single node)", "test_filter": "GinMPIDeviceTests.BarrierPools_NoCrossTalk", - "num_ranks": 4, "num_gpus": 4 + "name": "WindowRegister_AfterDevCommCreate", + "description": "Symmetric window register after ncclDevCommCreate (AICOMRCCL-2241)", + "test_filter": "GinMPIDeviceTests.WindowRegister_AfterDevCommCreate" }, - { "name": "SignalAdd_AndShadow", "description": "GIN device signal-add and shadow counter", "test_filter": "GinMPIDeviceTests.SignalAdd_AndShadow" }, - { "name": "SymPtr_PutAndPutValue", "description": "GIN symmetric-pointer put and put-value", "test_filter": "GinMPIDeviceTests.SymPtr_PutAndPutValue" }, - { "name": "Alltoall_PureReference", "description": "GIN pure-reference alltoall kernel", "test_filter": "GinMPIDeviceTests.Alltoall_PureReference" }, - { "name": "Alltoall_WorldGinBarrier", "description": "GIN alltoall on the world-team GIN barrier pool", "test_filter": "GinMPIDeviceTests.Alltoall_WorldGinBarrierReference" }, { - "name": "AlltoallHybrid_Reference", "description": "GIN hybrid (LSA + GIN) alltoall reference (2 ranks/node)", "test_filter": "GinMPIDeviceTests.AlltoallHybrid_Reference", - "num_ranks": 4, "num_gpus": 4 + "name": "SignalAdd_AndShadow", + "description": "GIN device signal-add and shadow counter", + "test_filter": "GinMPIDeviceTests.SignalAdd_AndShadow" }, { - "name": "Alltoall_SdmaInternal", "description": "GIN-SDMA alltoall via ncclAlltoAll (SDMA only, 8 ranks)", "test_filter": "GinMPIDeviceTests.Alltoall_SdmaInternal", - "num_ranks": 8, "num_nodes": 1, "num_gpus": 8 + "name": "Alltoall_PureReference", + "description": "GIN pure-reference alltoall kernel (registers after ncclDevCommCreate)", + "test_filter": "GinMPIDeviceTests.Alltoall_PureReference" }, - { "name": "MultiContext_AllFourRoute","description": "GIN multi-context, all four routes", "test_filter": "GinMPIDeviceTests.MultiContext_AllFourRoute" }, { - "name": "MultiContext_NonPowerOf2", "description": "GIN multi-context with non-power-of-2 context count (NCONNECTIONS=3, NCONTEXTS=3)", "test_filter": "GinMPIDeviceTests.MultiContext_NonPowerOf2", - "env_variables": { "NCCL_GIN_NCONNECTIONS": "3", "NCCL_GIN_NCONTEXTS": "3" } - }, - { "name": "LargeBuffer_Sweep", "description": "GIN large-buffer size sweep", "test_filter": "GinMPIDeviceTests.LargeBuffer_Sweep", "timeout": 300 }, - { - "name": "Disable_Error", "description": "Negative: devCommCreate must reject GIN when NCCL_GIN_ENABLE=0", "test_filter": "GinMPIDeviceTests.Disable_Error", - "env_variables": { "NCCL_GIN_ENABLE": "0" } - }, - { - "name": "Invalid_SignalPool", "description": "Negative: oversized signal pool rejected at devCommCreate", "test_filter": "GinMPIDeviceTests.Invalid_SignalPool", - "env_variables": { "NCCL_GIN_SIGNAL_POOL_SIZE": "0x40000000" } - }, - { - "name": "Invalid_CounterPool", "description": "Negative: oversized counter pool rejected at devCommCreate", "test_filter": "GinMPIDeviceTests.Invalid_CounterPool", - "env_variables": { "NCCL_GIN_COUNTER_POOL_SIZE": "0x40000000" } - }, - { "name": "Teardown_NoLeaks", "description": "GIN teardown leaves no leaks", "test_filter": "GinMPIDeviceTests.Teardown_NoLeaks" }, - { "name": "Init_Destroy_Stress", "description": "GIN init/destroy stress loop", "test_filter": "GinMPIDeviceTests.Init_Destroy_Stress", "timeout": 300 }, - { "name": "Alltoall_CrossNode", "description": "GIN cross-node alltoall (self-skips unless >=16 ranks)", "test_filter": "GinMPIDeviceTests.Alltoall_CrossNode" }, - { "name": "Properties_NLsaTeams", "description": "GIN properties: number of LSA teams", "test_filter": "GinMPIDeviceTests.Properties_NLsaTeams" }, - { "name": "MultiContext_Exclusive", "description": "GIN multi-context exclusive contexts", "test_filter": "GinMPIDeviceTests.MultiContext_Exclusive" }, - { "name": "VASignal_Put", "description": "GIN VA-signal put", "test_filter": "GinMPIDeviceTests.VASignal_Put" }, - { - "name": "RailConnection_Create", "description": "GIN rail connection create; rail-only mode (NCCL_CROSS_NIC=0)", "test_filter": "GinMPIDeviceTests.RailConnection_Create", - "env_variables": { "NCCL_CROSS_NIC": "0" } - }, - { - "name": "ReduceScatter_Symmetric", "description": "Symmetric-memory ReduceScatter sum (2 ranks/node)", "test_filter": "GinMPIDeviceTests.ReduceScatter_Symmetric", - "num_ranks": 4, "num_gpus": 4 - }, - { - "name": "ReduceScatter_Symmetric_Avg", "description": "Symmetric-memory ReduceScatter avg (2 ranks/node)", "test_filter": "GinMPIDeviceTests.ReduceScatter_Symmetric_Avg", - "num_ranks": 4, "num_gpus": 4 + "name": "AlltoallHybrid_Reference", + "description": "GIN hybrid (LSA + GIN) alltoall reference (2 ranks/node)", + "test_filter": "GinMPIDeviceTests.AlltoallHybrid_Reference", + "num_ranks": 4, + "num_gpus": 4 } ] - }, - "gin_devmpi_sdma": { - "extends": "gin_devmpi_common", - "description": "SDMA is a single-node backend: every GinMPIDeviceTests case runs on a single node (num_nodes=1). Cross-node cases (Put_CrossNode/Alltoall_CrossNode/ReduceScatter_Symmetric*) self-skip.", - "num_nodes": 1, - "num_gpus": 2, - "env_variables": { - "NCCL_GIN_TYPE": "6" - }, - "tests": [ - { "name": "BarrierSession_LsaOnly", "num_gpus": 4 }, - { "name": "BarrierSession_Hybrid", "num_gpus": 4 }, - { "name": "BarrierPools_NoCrossTalk", "num_gpus": 4 }, - { "name": "AlltoallHybrid_Reference", "num_gpus": 4 }, - { "name": "ReduceScatter_Symmetric", "num_gpus": 4 }, - { "name": "ReduceScatter_Symmetric_Avg", "num_gpus": 4 } - ] } }, "test_suites": [ { - "name": "GIN Device MPI Tests - SDMA (single node)", - "description": "Every GinMPIDeviceTests case, one mpirun per test, on the SDMA backend (single node)", + "name": "GIN Device MPI Tests - SDMA (AICOMRCCL-2241)", + "description": "Anvil SDMA post-DevCommCreate window-register subset, one mpirun per test", "config": "gin_devmpi_sdma", "enabled": true } From c8e45c2c406c25535ca8b97faa9316ec97f76a41 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Sun, 13 Sep 2026 13:57:17 -0400 Subject: [PATCH 09/17] fix(rccl): address Kapil review on AICOMRCCL-2241 unwind and tests Unregister GIN/RMA before unlinking memHead on symMemoryObtain failure, null ginHostWins/ginDevWins after fail_simple deregister, and add (void) on symMemoryUnregister GIN teardown. Strengthen WindowRegister_AfterDevCommCreate with a GIN put round-trip and drop the standalone gin_devmpi_sdma.json runner. Co-authored-by: Cursor --- projects/rccl/src/dev_runtime.cc | 6 +- .../rccl/test/transport/GinDeviceMPITests.cpp | 70 +++++++++++--- .../test_runner/configs/gin_devmpi_sdma.json | 96 ------------------- 3 files changed, 64 insertions(+), 108 deletions(-) delete mode 100644 projects/rccl/tools/scripts/test_runner/configs/gin_devmpi_sdma.json diff --git a/projects/rccl/src/dev_runtime.cc b/projects/rccl/src/dev_runtime.cc index 37ef893cb15b..d6749a3b5c2f 100644 --- a/projects/rccl/src/dev_runtime.cc +++ b/projects/rccl/src/dev_runtime.cc @@ -720,6 +720,10 @@ static ncclResult_t symMemoryRegisterGin(struct ncclComm* comm, struct ncclDevrM } free(mem->ginSegmentInfos); mem->ginSegmentInfos = nullptr; + for (int i = 0; i < NCCL_GIN_MAX_CONNECTIONS * NCCL_GIN_MAX_ACTIVE_BACKENDS; i++) { + mem->ginHostWins[i] = nullptr; + mem->ginDevWins[i] = nullptr; + } return ret; } @@ -735,7 +739,7 @@ static void symMemoryUnregister(struct ncclComm* comm, struct ncclDevrMemory* me struct ncclDevrState* devr = &comm->devrState; if (devr->ginEnabled && mem->ginSegmentInfos != nullptr) { for (int segment = 0; segment < mem->numGinSegments; segment++) { - ncclGinDeregister(comm, mem->ginSegmentInfos[segment].ginHostWins); + (void)ncclGinDeregister(comm, mem->ginSegmentInfos[segment].ginHostWins); } } if (devr->rmaProxyEnabled && mem->maxGlobalNumSegments == 1) { diff --git a/projects/rccl/test/transport/GinDeviceMPITests.cpp b/projects/rccl/test/transport/GinDeviceMPITests.cpp index 95e0636bd27d..b0927d62ace6 100644 --- a/projects/rccl/test/transport/GinDeviceMPITests.cpp +++ b/projects/rccl/test/transport/GinDeviceMPITests.cpp @@ -2230,35 +2230,83 @@ TEST_F(GinMPIDeviceTests, WindowRegister_AfterDevCommCreate) { if (auto reason = ginProxyTestSkipReason(); !reason.empty()) GTEST_SKIP() << reason; - if (!validateTestPrerequisites(/*min_processes=*/2, /*max_processes=*/8)) - GTEST_SKIP() << "Requires 2-8 ranks"; + if (!validateTestPrerequisites(/*min_processes=*/2, /*max_processes=*/2)) + GTEST_SKIP() << "Requires exactly 2 ranks"; ASSERT_EQ(ncclSuccess, createTestCommunicator()); - ncclComm_t comm = getActiveCommunicator(); + ncclComm_t comm = getActiveCommunicator(); + hipStream_t stream = getActiveStream(); + + int rank = -1, nRanks = -1; + ncclCommUserRank(comm, &rank); + ncclCommCount(comm, &nRanks); + ASSERT_EQ(2, nRanks); ncclDevCommRequirements reqs = defaultGinReqs(); - reqs.ginSignalCount = 1; + reqs.railGinBarrierCount = 1; + reqs.ginSignalCount = 1; ncclDevComm devComm{}; ASSERT_MPI_EQ(ncclSuccess, ncclDevCommCreate(comm, &reqs, &devComm)); auto devCommCleanup = makeScopeGuard([&]() { (void)ncclDevCommDestroy(comm, &devComm); }); - constexpr size_t kBytes = 64; - void* dBuf = nullptr; - ASSERT_MPI_EQ(ncclSuccess, ncclMemAlloc(&dBuf, kBytes)); + constexpr size_t kBufBytes = 64; + constexpr size_t kTransferBytes = 64; + constexpr ncclGinSignal_t kSigIdx = 0; + constexpr int kPeer = 1; + + void* dSrc = nullptr; + void* dDst = nullptr; + ASSERT_MPI_EQ(ncclSuccess, ncclMemAlloc(&dSrc, kBufBytes)); + ASSERT_MPI_EQ(ncclSuccess, ncclMemAlloc(&dDst, kBufBytes)); auto memCleanup = makeScopeGuard([&]() { - if (dBuf) (void)ncclMemFree(dBuf); + if (dSrc) (void)ncclMemFree(dSrc); + if (dDst) (void)ncclMemFree(dDst); }); - ncclWindow_t win = nullptr; + ncclWindow_t srcWin = nullptr, dstWin = nullptr; ASSERT_MPI_EQ(ncclSuccess, - ncclCommWindowRegister(comm, dBuf, kBytes, &win, NCCL_WIN_COLL_SYMMETRIC)); + ncclCommWindowRegister(comm, dSrc, kBufBytes, &srcWin, NCCL_WIN_COLL_SYMMETRIC)); + ASSERT_MPI_EQ(ncclSuccess, + ncclCommWindowRegister(comm, dDst, kBufBytes, &dstWin, NCCL_WIN_COLL_SYMMETRIC)); auto winCleanup = makeScopeGuard([&]() { - if (win) (void)ncclCommWindowDeregister(comm, win); + if (srcWin) (void)ncclCommWindowDeregister(comm, srcWin); + if (dstWin) (void)ncclCommWindowDeregister(comm, dstWin); }); + std::vector hostSrc(kBufBytes, 0); + std::vector hostDst(kBufBytes, 0); + for (size_t i = 0; i < kTransferBytes; i++) { + hostSrc[i] = static_cast(0xA0 + (i & 0x3F)); + } + ASSERT_MPI_EQ(hipSuccess, hipMemcpy(dSrc, hostSrc.data(), kBufBytes, hipMemcpyHostToDevice)); + ASSERT_MPI_EQ(hipSuccess, hipMemcpy(dDst, hostDst.data(), kBufBytes, hipMemcpyHostToDevice)); + + MPI_Barrier(MPI_COMM_WORLD); + + if (rank == 0) { + putBasicProducerKernel<<>>( + srcWin, /*srcOff=*/0, + dstWin, /*dstOff=*/0, + kTransferBytes, kSigIdx, kPeer, + devComm); + } else { + putBasicConsumerKernel<<>>( + kSigIdx, /*expectedSignalValue=*/1, devComm); + } + ASSERT_MPI_EQ(hipSuccess, hipStreamSynchronize(stream)); + MPI_Barrier(MPI_COMM_WORLD); + + if (rank == 1) { + std::vector hostResult(kBufBytes, 0); + ASSERT_EQ(hipSuccess, + hipMemcpy(hostResult.data(), dDst, kBufBytes, hipMemcpyDeviceToHost)); + for (size_t i = 0; i < kTransferBytes; i++) { + ASSERT_EQ(hostSrc[i], hostResult[i]) << "byte " << i << " mismatched"; + } + } } // Collective over the world team: composes the inner LSA barrier with the diff --git a/projects/rccl/tools/scripts/test_runner/configs/gin_devmpi_sdma.json b/projects/rccl/tools/scripts/test_runner/configs/gin_devmpi_sdma.json deleted file mode 100644 index 639ac99edab5..000000000000 --- a/projects/rccl/tools/scripts/test_runner/configs/gin_devmpi_sdma.json +++ /dev/null @@ -1,96 +0,0 @@ -{ - "system_configurations": { - "name": "RCCL-GIN-Device-MPI-SDMA-MI300X-Mellanox-IB", - "description": "Anvil SDMA (NCCL_GIN_TYPE=6) GinMPIDeviceTests that catch post-ncclDevCommCreate window register. Single node, no GDA, no alltoall_perf." - }, - "paths": { - "workdir": "${WORKDIR:-$PWD}", - "rocm_path": "${ROCM_PATH:-/opt/rocm}", - "mpi_path": "${MPI_PATH:-/opt/shared/ompi}" - }, - "env_variables": { - "NCCL_SOCKET_IFNAME": "eth0,eth1", - "HSA_NO_SCRATCH_RECLAIM": "1", - "NCCL_DEBUG": "INFO" - }, - "build_configuration": { - "install_flags": [ - "-l", - "--debug", - "-t", - "--enable-mpi-tests", - "--enable-code-coverage", - "--rocshmem-gin", - "--prefix", - "${WORKDIR}/install" - ], - "parallel_jobs": 64 - }, - "rccl_tests_build_configuration": { - "enabled": false - }, - "test_configurations": { - "default": { - "env_variables": {}, - "rerun_env_variables": { - "NCCL_DEBUG": "INFO" - } - }, - "gin_devmpi_sdma": { - "extends": "default", - "description": "AICOMRCCL-2241 Anvil SDMA subset: WindowRegister_AfterDevCommCreate plus the device-API cases that also register after ncclDevCommCreate. Omits cross-node, VA-signal, and Anvil ReduceScatter cases that self-skip at NCCL_GIN_TYPE=6 on one node.", - "is_gtest": true, - "binary": "rccl-UnitTestsMPI", - "num_ranks": 2, - "num_nodes": 1, - "num_gpus": 2, - "timeout": 180, - "env_variables": { - "NCCL_GIN_ENABLE": "1", - "NCCL_CUMEM_ENABLE": "1", - "NCCL_ENV_PLUGIN": "none", - "NCCL_DMABUF_ENABLE": "1", - "NCCL_GIN_TYPE": "6" - }, - "tests": [ - { - "name": "BarrierSession_LsaOnly", - "description": "ncclBarrierSession LSA-only (2 ranks/node, non-trivial LSA team)", - "test_filter": "GinMPIDeviceTests.BarrierSession_LsaOnly", - "num_ranks": 4, - "num_gpus": 4 - }, - { - "name": "WindowRegister_AfterDevCommCreate", - "description": "Symmetric window register after ncclDevCommCreate (AICOMRCCL-2241)", - "test_filter": "GinMPIDeviceTests.WindowRegister_AfterDevCommCreate" - }, - { - "name": "SignalAdd_AndShadow", - "description": "GIN device signal-add and shadow counter", - "test_filter": "GinMPIDeviceTests.SignalAdd_AndShadow" - }, - { - "name": "Alltoall_PureReference", - "description": "GIN pure-reference alltoall kernel (registers after ncclDevCommCreate)", - "test_filter": "GinMPIDeviceTests.Alltoall_PureReference" - }, - { - "name": "AlltoallHybrid_Reference", - "description": "GIN hybrid (LSA + GIN) alltoall reference (2 ranks/node)", - "test_filter": "GinMPIDeviceTests.AlltoallHybrid_Reference", - "num_ranks": 4, - "num_gpus": 4 - } - ] - } - }, - "test_suites": [ - { - "name": "GIN Device MPI Tests - SDMA (AICOMRCCL-2241)", - "description": "Anvil SDMA post-DevCommCreate window-register subset, one mpirun per test", - "config": "gin_devmpi_sdma", - "enabled": true - } - ] -} From 791bc629739ab1afdc4bdfb169f4c58ce0df82b9 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Wed, 16 Sep 2026 21:48:20 -0400 Subject: [PATCH 10/17] fix(rccl): unmap LSA slices on symMemoryObtain failure (AICOMRCCL-2241) fail_mem_space freed the big-space offset while per-rank cuMemMap slices were still live. Share the destroy unmap helper and cover it from the GIN-register failure host test. --- projects/rccl/src/dev_runtime.cc | 33 ++++++++++++++------- projects/rccl/test/DevRuntimeTestsStubs.cc | 12 ++++++-- projects/rccl/test/gin_teardown_test.cpp | 34 ++++++++++++++++++++++ 3 files changed, 66 insertions(+), 13 deletions(-) diff --git a/projects/rccl/src/dev_runtime.cc b/projects/rccl/src/dev_runtime.cc index d6749a3b5c2f..cfe93de7792e 100644 --- a/projects/rccl/src/dev_runtime.cc +++ b/projects/rccl/src/dev_runtime.cc @@ -747,6 +747,26 @@ static void symMemoryUnregister(struct ncclComm* comm, struct ncclDevrMemory* me } } +// Unmap LSA-flat slices created by symMemoryMapLsaTeam. Must run before +// ncclSpaceFree so a later obtain can remap the same bigOffset, and so +// ncclDevrFinalize's cuMemAddressFree does not see leftover mappings. +static void symMemoryUnmapLsaTeam(struct ncclComm* comm, struct ncclDevrMemory* mem) { + if (comm == nullptr || mem == nullptr || mem->lsaNumSegments == nullptr) return; + struct ncclDevrState* devr = &comm->devrState; + if (devr->lsaFlatBase == nullptr) return; + for (int r = 0; r < devr->lsaSize; r++) { + uintptr_t base = reinterpret_cast(devr->lsaFlatBase); + uintptr_t addr = base + r * devr->bigSize + mem->bigOffset; + for (int idx = 0; idx < mem->lsaNumSegments[r]; idx++) { + CUdeviceptr tmpBase; + size_t tmpBaseSize; + CUCHECKIGNORE(cuMemGetAddressRange(&tmpBase, &tmpBaseSize, reinterpret_cast(addr))); + CUCHECKIGNORE(cuMemUnmap(reinterpret_cast(addr), tmpBaseSize)); + addr = addr + tmpBaseSize; + } + } +} + // On success we take caller's reference on memHandle. // Due to multicast binds for each pre-exiting team, this function requires // caller do a world barrier before returning to user. @@ -909,6 +929,7 @@ static ncclResult_t symMemoryObtain(struct ncclComm* comm, CUmemGenericAllocatio symUnbindTeamLe(comm, mem, t->mcLeId); } fail_mem_space: + symMemoryUnmapLsaTeam(comm, mem); ncclSpaceFree(&devr->bigSpace, bigOffset, mem->lsaMaxSize); fail_mem: if (mem != nullptr) { @@ -944,17 +965,7 @@ static void symMemoryDestroy(struct ncclComm* comm, struct ncclDevrMemory* mem) symUnbindTeamMemory(comm, t, mem); symUnbindTeamLe(comm, mem, t->mcLeId); } - for (int r = 0; r < devr->lsaSize; r++) { - uintptr_t base = reinterpret_cast(devr->lsaFlatBase); - uintptr_t addr = base + r * devr->bigSize + mem->bigOffset; - for (int idx = 0; idx < mem->lsaNumSegments[r]; idx++) { - CUdeviceptr tmpBase; - size_t tmpBaseSize; - CUCHECKIGNORE(cuMemGetAddressRange(&tmpBase, &tmpBaseSize, reinterpret_cast(addr))); - CUCHECKIGNORE(cuMemUnmap(reinterpret_cast(addr), tmpBaseSize)); - addr = addr + tmpBaseSize; - } - } + symMemoryUnmapLsaTeam(comm, mem); ncclSpaceFree(&devr->bigSpace, mem->bigOffset, mem->lsaMaxSize); for (int segment = 0; segment < mem->numSegments; segment++) { diff --git a/projects/rccl/test/DevRuntimeTestsStubs.cc b/projects/rccl/test/DevRuntimeTestsStubs.cc index 9071a84640ab..644f445afd30 100644 --- a/projects/rccl/test/DevRuntimeTestsStubs.cc +++ b/projects/rccl/test/DevRuntimeTestsStubs.cc @@ -44,6 +44,8 @@ // Count hipMemAddressFree for skip-on vs skip-off finalize tests. int rcclTestHipMemAddressFreeCount = 0; +int rcclTestHipMemMapCount = 0; +int rcclTestHipMemUnmapCount = 0; // --------------------------------------------------------------------------- // Globals the translation unit references. @@ -355,10 +357,14 @@ HIP_FAKE hipError_t hipMemImportFromShareableHandle(hipMemGenericAllocationHandl return hipSuccess; } HIP_FAKE hipError_t hipMemMap(void*, size_t, size_t, hipMemGenericAllocationHandle_t, unsigned long long) { + rcclTestHipMemMapCount++; return hipSuccess; } HIP_FAKE hipError_t hipMemSetAccess(void*, size_t, const hipMemAccessDesc*, size_t) { return hipSuccess; } -HIP_FAKE hipError_t hipMemUnmap(void*, size_t) { return hipSuccess; } +HIP_FAKE hipError_t hipMemUnmap(void*, size_t) { + rcclTestHipMemUnmapCount++; + return hipSuccess; +} HIP_FAKE hipError_t hipMemRelease(hipMemGenericAllocationHandle_t) { return hipSuccess; } HIP_FAKE hipError_t hipMemRetainAllocationHandle(hipMemGenericAllocationHandle_t* handle, void*) { if (handle) *handle = reinterpret_cast(0x1); @@ -366,7 +372,9 @@ HIP_FAKE hipError_t hipMemRetainAllocationHandle(hipMemGenericAllocationHandle_t } HIP_FAKE hipError_t hipMemGetAddressRange(hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr) { if (pbase) *pbase = dptr; - if (psize) *psize = 0; + // Host tests map one 4096-byte segment per LSA rank. Returning 0 would still + // call unmap once (the idx loop advances), but a real size matches destroy. + if (psize) *psize = 4096; return hipSuccess; } diff --git a/projects/rccl/test/gin_teardown_test.cpp b/projects/rccl/test/gin_teardown_test.cpp index 2555fc353205..313e7f78d4eb 100644 --- a/projects/rccl/test/gin_teardown_test.cpp +++ b/projects/rccl/test/gin_teardown_test.cpp @@ -19,6 +19,8 @@ #include "common/ProcessIsolatedTestRunner.hpp" extern int rcclTestHipMemAddressFreeCount; +extern int rcclTestHipMemMapCount; +extern int rcclTestHipMemUnmapCount; extern "C" void DevRuntimeTests_SetGinRegisterFail(int fail); extern "C" struct ncclDevrMemory* DevRuntimeTests_GinRegisterMemHeadAtCall(); @@ -131,6 +133,38 @@ TEST_F(SymMemoryObtainTest, ObtainFailureUnlinksMemHeadWhenGinRegisterFails) { DevRuntimeTests_SetGinRegisterFail(0); } +// fail_mem_space_teams must unmap LSA-flat slices before ncclSpaceFree, matching +// symMemoryDestroy. Otherwise a later obtain remaps the same offset while the +// previous VA is still mapped. +TEST_F(SymMemoryObtainTest, ObtainFailureUnmapsLsaSlicesWhenGinRegisterFails) { + comm->devrState.ginEnabled = true; + DevRuntimeTests_SetGinRegisterFail(1); + rcclTestHipMemMapCount = 0; + rcclTestHipMemUnmapCount = 0; + + hipMemGenericAllocationHandle_t memHandle = reinterpret_cast(0x1); + void* userAddr = reinterpret_cast(0x100000); + struct ncclDevrMemory* mem = nullptr; + + EXPECT_NE(symMemoryObtain(comm, &memHandle, /*numSegments=*/1, userAddr, /*size=*/4096, /*winFlags=*/0, &mem), + ncclSuccess); + EXPECT_GT(rcclTestHipMemMapCount, 0); + EXPECT_GE(rcclTestHipMemUnmapCount, 1); + EXPECT_EQ(comm->devrState.memHead, nullptr); + + DevRuntimeTests_SetGinRegisterFail(0); + rcclTestHipMemMapCount = 0; + rcclTestHipMemUnmapCount = 0; + + ASSERT_EQ(symMemoryObtain(comm, &memHandle, /*numSegments=*/1, userAddr, /*size=*/4096, /*winFlags=*/0, &mem), + ncclSuccess); + ASSERT_NE(mem, nullptr); + EXPECT_GT(rcclTestHipMemMapCount, 0); + symMemoryDestroy(comm, mem); + EXPECT_GE(rcclTestHipMemUnmapCount, 1); + EXPECT_EQ(comm->devrState.memHead, nullptr); +} + // --------------------------------------------------------------------------- // AICOMRCCL-835 finalize-drain coverage. // From f6873328a3365327a89adf0e015ce04bf7530849 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Wed, 16 Sep 2026 21:52:48 -0400 Subject: [PATCH 11/17] test(rccl): fill ginSegmentInfos in host stubs after GIN register split develop now builds segment infos before ncclGinRegister. The no-op stub left numGinSegments at 0, so the GIN-failure host tests never reached registration. --- projects/rccl/test/DevRuntimeTestsStubs.cc | 16 ++++++++++++++-- 1 file changed, 14 insertions(+), 2 deletions(-) diff --git a/projects/rccl/test/DevRuntimeTestsStubs.cc b/projects/rccl/test/DevRuntimeTestsStubs.cc index 644f445afd30..5174baf9ca8a 100644 --- a/projects/rccl/test/DevRuntimeTestsStubs.cc +++ b/projects/rccl/test/DevRuntimeTestsStubs.cc @@ -234,14 +234,26 @@ ncclResult_t ncclRmaProxyDeregister(struct ncclComm*, void*[NCCL_GIN_MAX_CONNECT // --------------------------------------------------------------------------- // devr internal helpers (defined elsewhere in the real build). // --------------------------------------------------------------------------- -ncclResult_t ncclDevrPopulateSegmentSizes(struct ncclDevrMemory*, int) { return ncclSuccess; } +ncclResult_t ncclDevrPopulateSegmentSizes(struct ncclDevrMemory* mem, int numSegments) { + if (mem != nullptr && mem->segmentSizes != nullptr && numSegments > 0) { + mem->segmentSizes[0] = mem->size; + } + return ncclSuccess; +} ncclResult_t ncclDevrAllocAndPopulateSegmentWindows(struct ncclDevrState*, struct ncclDevrMemory*, hipStream_t, struct ncclSegmentWindow** out) { if (out) *out = nullptr; return ncclSuccess; } ncclResult_t ncclDevrVerifySegmentLayouts(struct ncclDevrMemory*, struct ncclComm*) { return ncclSuccess; } -ncclResult_t ncclDevrBuildGinSegmentInfos(struct ncclDevrMemory*) { return ncclSuccess; } +ncclResult_t ncclDevrBuildGinSegmentInfos(struct ncclDevrMemory* mem) { + if (mem == nullptr) return ncclInternalError; + mem->numGinSegments = 1; + NCCLCHECK(ncclCalloc(&mem->ginSegmentInfos, 1)); + mem->ginSegmentInfos[0].segmentSize = mem->size; + mem->ginSegmentInfos[0].memType = CU_MEM_LOCATION_TYPE_DEVICE; + return ncclSuccess; +} // --------------------------------------------------------------------------- // CFT / LE helpers pulled in via #include of hipified dev_runtime.cc. From 300628a67a7fffe421ab5f3cae11eea0f191b7b7 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Thu, 17 Sep 2026 02:51:39 -0400 Subject: [PATCH 12/17] fix(rccl): address Argus unwind and host RAS stub build (AICOMRCCL-2241) Deregister the failing GIN segment on partial ncclGinRegister rollback and zero-init LSA unmap sizes when cuMemGetAddressRange fails. Drop ras_param.h from the host micro-test stub so CPU-only CI can compile the RAS symbols. Co-authored-by: Cursor --- projects/rccl/src/dev_runtime.cc | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/projects/rccl/src/dev_runtime.cc b/projects/rccl/src/dev_runtime.cc index cfe93de7792e..fe21a189968a 100644 --- a/projects/rccl/src/dev_runtime.cc +++ b/projects/rccl/src/dev_runtime.cc @@ -715,7 +715,7 @@ static ncclResult_t symMemoryRegisterGin(struct ncclComm* comm, struct ncclDevrM } return ret; fail: - for (int i = 0; i < numSegmentsRegistered; i++) { + for (int i = 0; mem->ginSegmentInfos != nullptr && i <= numSegmentsRegistered && i < mem->numGinSegments; i++) { (void)ncclGinDeregister(comm, mem->ginSegmentInfos[i].ginHostWins); } free(mem->ginSegmentInfos); @@ -758,8 +758,8 @@ static void symMemoryUnmapLsaTeam(struct ncclComm* comm, struct ncclDevrMemory* uintptr_t base = reinterpret_cast(devr->lsaFlatBase); uintptr_t addr = base + r * devr->bigSize + mem->bigOffset; for (int idx = 0; idx < mem->lsaNumSegments[r]; idx++) { - CUdeviceptr tmpBase; - size_t tmpBaseSize; + CUdeviceptr tmpBase = nullptr; + size_t tmpBaseSize = 0; CUCHECKIGNORE(cuMemGetAddressRange(&tmpBase, &tmpBaseSize, reinterpret_cast(addr))); CUCHECKIGNORE(cuMemUnmap(reinterpret_cast(addr), tmpBaseSize)); addr = addr + tmpBaseSize; From c6ec7942b3790e6bc805f47776eb90212c33575f Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Thu, 17 Sep 2026 15:45:22 -0400 Subject: [PATCH 13/17] test(rccl): drop RAS host stub; expect both GIN segments deregistered develop now supplies ras_param_fakes.cc, so the extra ras_fakes.cc is redundant. The widened GIN rollback also deregisters the part-way failing segment. Co-authored-by: Cursor --- projects/rccl/test/host/dev-runtime-test.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/rccl/test/host/dev-runtime-test.cc b/projects/rccl/test/host/dev-runtime-test.cc index 501f3928d199..876cfb24c273 100644 --- a/projects/rccl/test/host/dev-runtime-test.cc +++ b/projects/rccl/test/host/dev-runtime-test.cc @@ -1956,7 +1956,7 @@ TEST_F(SymMemoryRegisterGinElasticTest, SecondSegmentFails_DeregistersTheFirst) EXPECT_NE(symMemoryRegisterGin(comm, &mem), ncclSuccess); EXPECT_EQ(reg.calls, 2); - EXPECT_EQ(dereg.calls, 1); // exactly the one that succeeded + EXPECT_EQ(dereg.calls, 2); // first segment plus the part-way failing one EXPECT_EQ(mem.ginSegmentInfos, nullptr); } From 10ed86bca11af0dd9a089b98b89e878af7adb0f2 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Thu, 17 Sep 2026 18:44:02 -0400 Subject: [PATCH 14/17] test(rccl): run WindowRegister_AfterDevCommCreate on gin_devmpi_sdma The Anvil SDMA subset replaces the common GinMPIDeviceTests list, so the AICOMRCCL-2241 case never ran on NCCL_GIN_TYPE=7. Keep the GDA common entry and note the changelog type after the 2.31 renumber. Co-authored-by: Cursor --- projects/rccl/CHANGELOG.md | 2 +- .../test_runner/configs/gin_rocshmem_backends_mellanox.json | 1 + 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md index 4b67b60f0e28..6c494f3f77cd 100644 --- a/projects/rccl/CHANGELOG.md +++ b/projects/rccl/CHANGELOG.md @@ -63,7 +63,7 @@ Full documentation for RCCL is available at [https://rccl.readthedocs.io](https: * Fixed gfx1250 LL and LL128 comm-FIFO hangs on sibling DPX partitions. The FIFO store now uses system-scope b128 (`RCCL_LL_FIFO_SYS_SCOPE`) alongside the existing system-scope load, preventing hangs at slot reuse starting from the 9th collective operation. * Fixed DDA fabric AllToAll validation race by staging send data into scratch with a host-launched `cudaMemcpyAsync` before the peer exchange kernel. * Fixed DDA fabric barrier publication race causing sporadic validation errors by using release-acquire semantics on the prologue barrier. -* Fixed `ncclCommWindowRegister` after `ncclDevCommCreate` aborting on the GIN Anvil SDMA backend (`NCCL_GIN_TYPE=6`) with `could not resolve LSA flat addr`. Symmetric-window memory is now linked onto the device-runtime list before GIN/RMA registration so plugins that resolve the user VA through that list can see the in-flight window. +* Fixed `ncclCommWindowRegister` after `ncclDevCommCreate` aborting on the GIN Anvil SDMA backend (`NCCL_GIN_TYPE=7`) with `could not resolve LSA flat addr`. Symmetric-window memory is now linked onto the device-runtime list before GIN/RMA registration so plugins that resolve the user VA through that list can see the in-flight window. ### Known issues * On gfx90a (MI210/MI250/MI250X) with ROCm 7.13 or later, per-launch scratch-memory reclaim in the runtime degrades RCCL performance. Set `HSA_NO_SCRATCH_RECLAIM=1` to restore performance. diff --git a/projects/rccl/tools/scripts/test_runner/configs/gin_rocshmem_backends_mellanox.json b/projects/rccl/tools/scripts/test_runner/configs/gin_rocshmem_backends_mellanox.json index 1b9f0b2653c8..237671b86974 100644 --- a/projects/rccl/tools/scripts/test_runner/configs/gin_rocshmem_backends_mellanox.json +++ b/projects/rccl/tools/scripts/test_runner/configs/gin_rocshmem_backends_mellanox.json @@ -212,6 +212,7 @@ }, "tests": [ { "name": "BarrierSession_LsaOnly", "num_gpus": 4 }, + { "name": "WindowRegister_AfterDevCommCreate" }, { "name": "BarrierSession_Hybrid", "num_gpus": 4 }, { "name": "BarrierPools_NoCrossTalk", "num_gpus": 4 }, { "name": "AlltoallHybrid_Reference", "num_gpus": 4 }, From 19dcaca0c0f1970b5bb16e4794fe58494d26b5fb Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Fri, 18 Sep 2026 15:47:56 -0400 Subject: [PATCH 15/17] docs(rccl): sync symMemoryObtain stage order with memHead-before-GIN fix Swap the design-doc stages and tighten comments from Argus/Aurelien review on AICOMRCCL-2241 so the in-tree description matches link-before-register. Co-authored-by: Cursor --- .../NCCL_Gin_and_Symmetric_Memory.md | 5 +++-- projects/rccl/src/dev_runtime.cc | 8 ++++---- projects/rccl/src/gin/gin_host.cc | 3 --- projects/rccl/src/rma/rma_proxy.cc | 3 --- 4 files changed, 7 insertions(+), 12 deletions(-) diff --git a/projects/rccl/docs/contrib/GIN/NCCL_Gin_and_Symmetric_Memory/NCCL_Gin_and_Symmetric_Memory.md b/projects/rccl/docs/contrib/GIN/NCCL_Gin_and_Symmetric_Memory/NCCL_Gin_and_Symmetric_Memory.md index 267bb068b177..56ba0e134898 100644 --- a/projects/rccl/docs/contrib/GIN/NCCL_Gin_and_Symmetric_Memory/NCCL_Gin_and_Symmetric_Memory.md +++ b/projects/rccl/docs/contrib/GIN/NCCL_Gin_and_Symmetric_Memory/NCCL_Gin_and_Symmetric_Memory.md @@ -243,8 +243,9 @@ symMemoryObtain ├── Stage 5: symMemoryMapLsaTeam constructs lsaFlatBase VA + maps each peer's physical segment ├── Stage 6: backfill primaryAddr (for callers that didn't supply a VA) ├── Stage 7: symBindTeamMemory binds NVLS multicast - ├── Stage 8: symMemoryRegisterGin / symMemoryRegisterRma — cross-node registration - └── Stage 9: link into devr->memHead, refCount=1, return + ├── Stage 8: link into devr->memHead (plugins resolve user VA via memHead) + ├── Stage 9: symMemoryRegisterGin / symMemoryRegisterRma — cross-node registration + └── return with refCount=1 ``` On NVLink, the local rank's handle is exported as follows: diff --git a/projects/rccl/src/dev_runtime.cc b/projects/rccl/src/dev_runtime.cc index fe21a189968a..0d6769314d44 100644 --- a/projects/rccl/src/dev_runtime.cc +++ b/projects/rccl/src/dev_runtime.cc @@ -886,10 +886,10 @@ static ncclResult_t symMemoryObtain(struct ncclComm* comm, CUmemGenericAllocatio } } - // Link before GIN/RMA registration. Plugins such as Anvil SDMA resolve the - // user VA through ncclDevrGetLsaSelfAddr, which only searches memHead (and - // the LSA flat range). Registering after ncclDevCommCreate is legal, so the - // in-flight mem must already be on the list. Unlink on the failure path. + // Add to the list of mems before GIN/RMA registration. Plugins such as Anvil SDMA resolve the + // user VA through ncclDevrGetLsaSelfAddr, which only searches memHead (and the LSA flat range). + // Registering after ncclDevCommCreate is legal, so the registered mem must be on the list before + // GIN assignment. Unlink on the failure path. mem->next = devr->memHead; devr->memHead = mem; diff --git a/projects/rccl/src/gin/gin_host.cc b/projects/rccl/src/gin/gin_host.cc index a1fb1990fbfe..093034846827 100644 --- a/projects/rccl/src/gin/gin_host.cc +++ b/projects/rccl/src/gin/gin_host.cc @@ -547,9 +547,6 @@ ncclResult_t ncclGinDeregister(struct ncclComm* comm, struct ncclGinBackendState* backend = &ginState->backends[backendIdx]; for (int commIdx = 0; commIdx < backend->ginCommCount; commIdx++) { int slot = backendIdx * NCCL_GIN_MAX_CONNECTIONS + commIdx; - // ncclGinRegister returns at the first failing index, and the multi-segment - // DMABUF bail leaves every slot null. deregMrSym implementations deref the - // handle without a null check. if (ginHostWins[slot] == nullptr) continue; NCCLCHECK(backend->ncclGin->deregMrSym(backend->ginComms[commIdx], ginHostWins[slot])); } diff --git a/projects/rccl/src/rma/rma_proxy.cc b/projects/rccl/src/rma/rma_proxy.cc index f1a389a61af7..8d9774a1adf0 100644 --- a/projects/rccl/src/rma/rma_proxy.cc +++ b/projects/rccl/src/rma/rma_proxy.cc @@ -410,9 +410,6 @@ ncclResult_t ncclRmaProxyRegister(struct ncclComm* comm, void* address, size_t s ncclResult_t ncclRmaProxyDeregister(struct ncclComm* comm, void* rmaHostWins[NCCL_RMA_MAX_CONNECTIONS]) { struct ncclRmaProxyState* rmaProxyState = &comm->rmaState.rmaProxyState; for (int n = 0; n < rmaProxyState->rmaCommCount; n++) { - // ncclRmaProxyRegister returns at the first failing index, so later slots - // (and, on a partial success, any unused connections) stay null. deregMrSym - // implementations deref the handle without a null check. if (rmaHostWins[n] == nullptr) continue; NCCLCHECK(rmaProxyState->ncclRma->deregMrSym(rmaProxyState->rmaComms[n], rmaHostWins[n])); } From cbb71f33b7ba6af2838a25fa95cdf5fe9f621e35 Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Fri, 18 Sep 2026 16:27:01 -0400 Subject: [PATCH 16/17] test(rccl): use HIP mem location enum in DevRuntime stubs (AICOMRCCL-2241) Math-CI ROCm 7.0.2 backwards-compat builds the host DevRuntime stubs without the CUDA-named CU_MEM_* aliases from p2p.h, so CU_MEM_LOCATION_TYPE_DEVICE failed to compile. Match sibling tests that already use hipMemLocationTypeDevice. Co-authored-by: Cursor --- projects/rccl/test/DevRuntimeTestsStubs.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/rccl/test/DevRuntimeTestsStubs.cc b/projects/rccl/test/DevRuntimeTestsStubs.cc index 5174baf9ca8a..eae8610dbe04 100644 --- a/projects/rccl/test/DevRuntimeTestsStubs.cc +++ b/projects/rccl/test/DevRuntimeTestsStubs.cc @@ -251,7 +251,7 @@ ncclResult_t ncclDevrBuildGinSegmentInfos(struct ncclDevrMemory* mem) { mem->numGinSegments = 1; NCCLCHECK(ncclCalloc(&mem->ginSegmentInfos, 1)); mem->ginSegmentInfos[0].segmentSize = mem->size; - mem->ginSegmentInfos[0].memType = CU_MEM_LOCATION_TYPE_DEVICE; + mem->ginSegmentInfos[0].memType = hipMemLocationTypeDevice; return ncclSuccess; } From deddfd027da65e097fc5df1d26b050d423ae72ad Mon Sep 17 00:00:00 2001 From: Donglai Dai Date: Tue, 22 Sep 2026 16:31:37 -0400 Subject: [PATCH 17/17] fix(rccl): guard symMemoryUnregister RMA deregister on rmaHostWins[0] Match the other deregister sites so a partway ncclRmaProxyConnectOnce failure cannot walk past the rmaHostWins slots. Note the ginWins clear loop as defensive-only per Argus; unlink helper and tests stay in AICOMRCCL-2428. Co-authored-by: Cursor --- projects/rccl/src/dev_runtime.cc | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/projects/rccl/src/dev_runtime.cc b/projects/rccl/src/dev_runtime.cc index 0d6769314d44..65fa6d049da0 100644 --- a/projects/rccl/src/dev_runtime.cc +++ b/projects/rccl/src/dev_runtime.cc @@ -720,6 +720,9 @@ static ncclResult_t symMemoryRegisterGin(struct ncclComm* comm, struct ncclDevrM } free(mem->ginSegmentInfos); mem->ginSegmentInfos = nullptr; + // Defensive only: mem-level ginHostWins/ginDevWins are populated solely by + // the numGinSegments==1 cache copy above, which sits after the last goto fail. + // ncclGinRegister writes ginSegmentInfos[seg].ginHostWins instead. for (int i = 0; i < NCCL_GIN_MAX_CONNECTIONS * NCCL_GIN_MAX_ACTIVE_BACKENDS; i++) { mem->ginHostWins[i] = nullptr; mem->ginDevWins[i] = nullptr; @@ -742,7 +745,11 @@ static void symMemoryUnregister(struct ncclComm* comm, struct ncclDevrMemory* me (void)ncclGinDeregister(comm, mem->ginSegmentInfos[segment].ginHostWins); } } - if (devr->rmaProxyEnabled && mem->maxGlobalNumSegments == 1) { + // rmaHostWins[0] is a reliable witness that register completed (same pattern + // as windowRegisterNonSym / ncclCommWindowDeregister). Skip if connect/register + // failed partway so deregister does not walk an unbound rmaCommCount. + if (devr->rmaProxyEnabled && mem->maxGlobalNumSegments == 1 && + mem->rmaHostWins[0] != nullptr) { (void)ncclRmaProxyDeregister(comm, mem->rmaHostWins); } }