dlamd1dai · Sep 24, 2026 · Sep 1, 2026 · Sep 4, 2026 · Sep 11, 2026 · Sep 11, 2026
@@ -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.

@@ -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:

@@ -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<uintptr_t>(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<CUdeviceptr>(addr)));
+      CUCHECKIGNORE(cuMemUnmap(reinterpret_cast<CUdeviceptr>(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<uintptr_t>(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<CUdeviceptr>(addr)));
-        CUCHECKIGNORE(cuMemUnmap(reinterpret_cast<CUdeviceptr>(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) {

@@ -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]));
     }
   }

@@ -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;

@@ -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,18 +369,24 @@ 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<hipMemGenericAllocationHandle_t>(0x1);
   return hipSuccess;
 }
 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;
 }