diff --git a/projects/rccl/CHANGELOG.md b/projects/rccl/CHANGELOG.md index ad9a3f5164dd..6c494f3f77cd 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=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/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 0cd9f6e4eec3..65fa6d049da0 100644 --- a/projects/rccl/src/dev_runtime.cc +++ b/projects/rccl/src/dev_runtime.cc @@ -713,15 +713,21 @@ 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); + for (int i = 0; mem->ginSegmentInfos != nullptr && i <= numSegmentsRegistered && i < mem->numGinSegments; i++) { + (void)ncclGinDeregister(comm, mem->ginSegmentInfos[i].ginHostWins); } free(mem->ginSegmentInfos); mem->ginSegmentInfos = nullptr; - goto exit; + // 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; + } + return ret; } static ncclResult_t symMemoryRegisterRma(struct ncclComm* comm, struct ncclDevrMemory* mem) { @@ -732,6 +738,42 @@ 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++) { + (void)ncclGinDeregister(comm, mem->ginSegmentInfos[segment].ginHostWins); + } + } + // 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); + } +} + +// 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 = nullptr; + size_t tmpBaseSize = 0; + 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. @@ -851,6 +893,13 @@ static ncclResult_t symMemoryObtain(struct ncclComm* comm, CUmemGenericAllocatio } } + // 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; + if (devr->ginEnabled) { NCCLCHECKGOTO(symMemoryRegisterGin(comm, mem), ret, fail_mem_space_teams); } else { @@ -866,15 +915,17 @@ 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: + symMemoryUnregister(comm, mem); + { + struct ncclDevrMemory** ptr = &devr->memHead; + while (*ptr != nullptr && *ptr != mem) ptr = &(*ptr)->next; + if (*ptr == mem) *ptr = mem->next; + } for (struct ncclDevrTeam* t = devr->teamHead; t != nullptr; t = t->next) { symUnbindTeamMemory(comm, t, mem); if (ucBound && t->ucLeId != NCCL_LE_ID_INVALID) { @@ -885,9 +936,11 @@ 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) { + free(mem->ginSegmentInfos); free(mem->memHandles); free(mem->segmentSizes); free(mem->lsaNumSegments); @@ -899,59 +952,40 @@ 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; - } - } - - ncclSpaceFree(&devr->bigSpace, mem->bigOffset, mem->lsaMaxSize); - for (int segment = 0; segment < mem->numSegments; segment++) { - CUCHECKIGNORE(cuMemRelease(mem->memHandles[segment])); - } + 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; + } - struct ncclDevrMemory** ptr = &devr->memHead; - while (*ptr != mem) ptr = &(*ptr)->next; - *ptr = mem->next; // Remove from list. + 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); + } + symMemoryUnmapLsaTeam(comm, mem); - 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/src/gin/gin_host.cc b/projects/rccl/src/gin/gin_host.cc index ac0ef27cfc76..093034846827 100644 --- a/projects/rccl/src/gin/gin_host.cc +++ b/projects/rccl/src/gin/gin_host.cc @@ -547,6 +547,7 @@ 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; + 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 8d418aa6d968..8d9774a1adf0 100644 --- a/projects/rccl/src/rma/rma_proxy.cc +++ b/projects/rccl/src/rma/rma_proxy.cc @@ -410,6 +410,7 @@ 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++) { + if (rmaHostWins[n] == nullptr) continue; NCCLCHECK(rmaProxyState->ncclRma->deregMrSym(rmaProxyState->rmaComms[n], rmaHostWins[n])); } return ncclSuccess; diff --git a/projects/rccl/test/DevRuntimeTestsStubs.cc b/projects/rccl/test/DevRuntimeTestsStubs.cc index bf6777ec81e7..eae8610dbe04 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. @@ -185,6 +187,14 @@ 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; + 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; @@ -199,8 +209,10 @@ 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; } ncclResult_t ncclGinDeregister(struct ncclComm*, void*[NCCL_GIN_MAX_CONNECTIONS]) { return ncclSuccess; } @@ -222,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 = hipMemLocationTypeDevice; + return ncclSuccess; +} // --------------------------------------------------------------------------- // CFT / LE helpers pulled in via #include of hipified dev_runtime.cc. @@ -345,10 +369,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); @@ -356,7 +384,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 381969e99c0f..313e7f78d4eb 100644 --- a/projects/rccl/test/gin_teardown_test.cpp +++ b/projects/rccl/test/gin_teardown_test.cpp @@ -19,6 +19,10 @@ #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(); // Build the smallest ncclComm/ncclDevrState that symMemoryObtain will accept: // a single-rank, single-LSA-team comm with GIN and RMA proxy disabled. @@ -111,6 +115,56 @@ 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_NE(DevRuntimeTests_GinRegisterMemHeadAtCall(), nullptr); + EXPECT_EQ(comm->devrState.memHead, nullptr); + + 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. // 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); } diff --git a/projects/rccl/test/transport/GinDeviceMPITests.cpp b/projects/rccl/test/transport/GinDeviceMPITests.cpp index f25988ffba2e..b0927d62ace6 100644 --- a/projects/rccl/test/transport/GinDeviceMPITests.cpp +++ b/projects/rccl/test/transport/GinDeviceMPITests.cpp @@ -2221,6 +2221,94 @@ 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=*/2)) + GTEST_SKIP() << "Requires exactly 2 ranks"; + + ASSERT_EQ(ncclSuccess, createTestCommunicator()); + 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.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 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 (dSrc) (void)ncclMemFree(dSrc); + if (dDst) (void)ncclMemFree(dDst); + }); + + ncclWindow_t srcWin = nullptr, dstWin = nullptr; + ASSERT_MPI_EQ(ncclSuccess, + 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 (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 // 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..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 @@ -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 @@ -211,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 },