diff --git a/projects/rccl/src/gin/gin_plugin_anvil_sdma.cc b/projects/rccl/src/gin/gin_plugin_anvil_sdma.cc index a542c4e98ac8..e9a1a5046da9 100644 --- a/projects/rccl/src/gin/gin_plugin_anvil_sdma.cc +++ b/projects/rccl/src/gin/gin_plugin_anvil_sdma.cc @@ -29,8 +29,10 @@ #include #include #include +#include static std::map bufferRegRefcount; +static std::map g_signalSpanRefcount; static std::mutex pluginMutex; // Per-rank comm objects whose LSA-signal peer connectivity has already been @@ -55,22 +57,40 @@ struct ginAnvilCollCtx { struct ginAnvilGinCtx { ncclNetDeviceHandle_v11_t* devHandle; ncclGinAnvilSdmaGPUContext* gpuCtxDev; - ncclGinAnvilSdmaGPUContext gpuCtxHost; + ncclGinAnvilSdmaGPUContext* gpuCtxHost; struct ncclComm* comm; int nRanks; int rank; + int nContexts; int nSignals; int nCounters; - int signalSlot; bool hasError; - bool signalsBound; + uint64_t* counters; void** gpu_queue_handles; uint64_t* sdma_dirty_d; int numChannels; int sdmaChannelStride; - uintptr_t* signal_remote_addrs_dev; + uintptr_t** signal_remote_addrs_dev; + void* signalSpanLsaSelf; }; +static void ginAnvilSignalSpanRefInc(void* spanLsaSelf) { + if (!spanLsaSelf) return; + std::lock_guard lock(pluginMutex); + g_signalSpanRefcount[(uintptr_t)spanLsaSelf]++; +} + +static void ginAnvilSignalSpanRefDec(void* spanLsaSelf) { + if (!spanLsaSelf) return; + std::lock_guard lock(pluginMutex); + auto it = g_signalSpanRefcount.find((uintptr_t)spanLsaSelf); + if (it == g_signalSpanRefcount.end()) return; + if (--it->second <= 0) { + (void)ncclGinAnvilIpcTableUnregister(spanLsaSelf); + g_signalSpanRefcount.erase(it); + } +} + struct GinAnvilPendingEntry { ginAnvilGinCtx* ctx; GinAnvilPendingEntry* next; @@ -107,6 +127,24 @@ static void ginAnvilPendingClear(struct ncclComm* comm) { g_pendingByComm.erase(comm); } +static void ginAnvilFreeGinCtx(ginAnvilGinCtx* ctx) { + if (!ctx) return; + if (ctx->comm) ginAnvilPendingRemove(ctx->comm, ctx); + for (int contextId = 0; contextId < ctx->nContexts; contextId++) { + if (ctx->gpuCtxHost) ncclGinAnvilIpcTableUntrackContext(&ctx->gpuCtxHost[contextId]); + if (ctx->signal_remote_addrs_dev && ctx->signal_remote_addrs_dev[contextId]) { + CUDACHECKIGNORE(hipFree(ctx->signal_remote_addrs_dev[contextId])); + } + } + ginAnvilSignalSpanRefDec(ctx->signalSpanLsaSelf); + if (ctx->counters) CUDACHECKIGNORE(hipFree(ctx->counters)); + if (ctx->gpuCtxDev) CUDACHECKIGNORE(hipFree(ctx->gpuCtxDev)); + free(ctx->signal_remote_addrs_dev); + delete[] ctx->gpuCtxHost; + free(ctx->devHandle); + delete ctx; +} + struct ginAnvilMemHandle { ncclGinAnvilSdmaMemHandle* devHandle; void* addr; @@ -139,6 +177,7 @@ void ncclGinAnvilPluginTestResetHostState(void) { g_pendingByComm.erase(comm); } bufferRegRefcount.clear(); + g_signalSpanRefcount.clear(); ginAnvilConnCheckedComms.clear(); } @@ -445,7 +484,8 @@ static const char* ginAnvilConnCheckStepName(GinAnvilConnCheckStep step) { } } -static ncclResult_t ginAnvilCheckSignalConnectivity(ginAnvilGinCtx* ctx, void* lsaSelf) { +static ncclResult_t ginAnvilCheckSignalConnectivity(ginAnvilGinCtx* ctx, void* lsaSelf, + uintptr_t* signalRemoteAddrsDev) { struct ncclComm* comm = ctx->comm; struct ncclDevrState* devr = &comm->devrState; const int nRanks = ctx->nRanks; @@ -498,10 +538,9 @@ static ncclResult_t ginAnvilCheckSignalConnectivity(ginAnvilGinCtx* ctx, void* l ctx->nSignals < nRanks) { setupState = kSetupSkip; } - if (setupState == kSetupReady && - (ctx->signal_remote_addrs_dev == nullptr || lsaSelf == nullptr)) { + if (setupState == kSetupReady && (signalRemoteAddrsDev == nullptr || lsaSelf == nullptr)) { WARN("GIN anvil-sdma: conn-check setup has null address (rank %d, remote-addrs=%p, lsaSelf=%p)", - rank, (void*)ctx->signal_remote_addrs_dev, lsaSelf); + rank, (void*)signalRemoteAddrsDev, lsaSelf); setupState = kSetupFailed; } if (setupState == kSetupReady && @@ -609,7 +648,7 @@ static ncclResult_t ginAnvilCheckSignalConnectivity(ginAnvilGinCtx* ctx, void* l if (rank == injRank) { WARN("GIN anvil-sdma: [TEST] injecting connectivity fault on rank %d (skipping signal writes)", rank); - } else if (ginAnvilConnWrite(ctx->signal_remote_addrs_dev, nRanks, rank, stamp, connStream) != 0) { + } else if (ginAnvilConnWrite(signalRemoteAddrsDev, nRanks, rank, stamp, connStream) != 0) { failAt(kConnCheckWrite); } @@ -718,60 +757,94 @@ static ncclResult_t ginAnvilCheckSignalConnectivity(ginAnvilGinCtx* ctx, void* l return ret; } -static ncclResult_t ginAnvilRegisterLsaSignals(ginAnvilGinCtx* ctx, void* lsaSelf, size_t bytes) { - struct ncclDevrState* devr = &ctx->comm->devrState; - struct ncclComm* comm = ctx->comm; +static ncclResult_t ginAnvilRegisterSignalSpan(struct ncclComm* comm, void* spanLsaSelf, size_t spanBytes, int nRanks, + int ginRank) { + struct ncclDevrState* devr = &comm->devrState; const ptrdiff_t stride = (ptrdiff_t)devr->bigSize; - if (ctx->nRanks != devr->lsaSize) { - WARN("GIN anvil-sdma: signal nRanks=%d != devr->lsaSize=%d (rank %d)", ctx->nRanks, devr->lsaSize, ctx->rank); + if (nRanks != devr->lsaSize) { + WARN("GIN anvil-sdma: signal nRanks=%d != devr->lsaSize=%d (rank %d)", nRanks, devr->lsaSize, ginRank); } - if (ctx->rank != devr->lsaSelf) { - WARN("GIN anvil-sdma: ctx->rank=%d != devr->lsaSelf=%d (signal peer indexing may be wrong)", ctx->rank, + if (ginRank != devr->lsaSelf) { + WARN("GIN anvil-sdma: ctx->rank=%d != devr->lsaSelf=%d (signal peer indexing may be wrong)", ginRank, devr->lsaSelf); } - ctx->gpuCtxHost.signals = (uint64_t*)lsaSelf; - - uintptr_t* hostAddrs = (uintptr_t*)calloc((size_t)ctx->nRanks, sizeof(uintptr_t)); + // signal_remote_addrs is indexed by GIN-team peer, matching resolveRemotePeerVa's + // (peer - rsCtx->rank). lsaSelf is a different space under NCCL_GIN_CONNECTION_RAIL. + uintptr_t* hostAddrs = (uintptr_t*)calloc((size_t)nRanks, sizeof(uintptr_t)); if (!hostAddrs) return ncclSystemError; - for (int pe = 0; pe < ctx->nRanks; pe++) { - hostAddrs[pe] = (uintptr_t)lsaSelf + static_cast(pe - ctx->rank) * stride; + for (int pe = 0; pe < nRanks; pe++) { + hostAddrs[pe] = (uintptr_t)spanLsaSelf + static_cast(pe - ginRank) * stride; } - uintptr_t selfExpected = (uintptr_t)lsaSelf; - if (hostAddrs[ctx->rank] != selfExpected) { - WARN("GIN anvil-sdma: signal_remote_addrs[self=%d]=%#lx != signals=%#lx (lsaSelf=%#lx stride=%zd)", ctx->rank, - (unsigned long)hostAddrs[ctx->rank], (unsigned long)selfExpected, (unsigned long)lsaSelf, (long)stride); + int rc = ncclGinAnvilIpcTableRegisterExplicit(spanLsaSelf, hostAddrs, nRanks, spanBytes); + free(hostAddrs); + if (rc != 0) { + WARN("GIN anvil-sdma: LSA signal span IPC table register failed for %p size %zu", spanLsaSelf, spanBytes); + return ncclSystemError; } - uintptr_t* gatheredLocalBases = (uintptr_t*)calloc((size_t)ctx->nRanks, sizeof(uintptr_t)); + // Probe the span base once. Every logical context's slot is a constant offset + // from it, so running this per context would cost one bootstrap round trip per + // context inside ncclDevCommCreate. + // + // bootstrapAllGather writes one slot per world rank, so the buffer is sized and + // indexed in world space. nRanks/ginRank are the GIN team, which is smaller + // under NCCL_GIN_CONNECTION_RAIL. + const int worldRanks = comm->nRanks; + uintptr_t* gatheredLocalBases = + (worldRanks > 0 && comm->rank >= 0 && comm->rank < worldRanks) + ? (uintptr_t*)calloc((size_t)worldRanks, sizeof(uintptr_t)) + : nullptr; if (gatheredLocalBases) { - gatheredLocalBases[ctx->rank] = (uintptr_t)lsaSelf; + gatheredLocalBases[comm->rank] = (uintptr_t)spanLsaSelf; if (ginAnvilBootstrapAllgather(comm->bootstrap, gatheredLocalBases, sizeof(uintptr_t)) != 0) { WARN("GIN anvil-sdma: signal local-base allgather failed"); - } else if (gatheredLocalBases[ctx->rank] != (uintptr_t)lsaSelf) { - WARN("GIN anvil-sdma: signal local-base allgather mismatch rank=%d got=%#lx expect=%#lx", ctx->rank, - (unsigned long)gatheredLocalBases[ctx->rank], (unsigned long)lsaSelf); + } else if (gatheredLocalBases[comm->rank] != (uintptr_t)spanLsaSelf) { + WARN("GIN anvil-sdma: signal local-base allgather mismatch rank=%d got=%#lx expect=%#lx", comm->rank, + (unsigned long)gatheredLocalBases[comm->rank], (unsigned long)spanLsaSelf); } free(gatheredLocalBases); } + return ncclSuccess; +} - int rc = ncclGinAnvilIpcTableRegisterExplicit(lsaSelf, hostAddrs, ctx->nRanks, bytes); - if (rc != 0) { - WARN("GIN anvil-sdma: LSA signal IPC table register failed for %p size %zu", lsaSelf, bytes); - free(hostAddrs); +static ncclResult_t ginAnvilBindContextSignals(ginAnvilGinCtx* ctx, int contextId, int signalSlot, void* lsaSelf, + size_t bytes) { + struct ncclDevrState* devr = &ctx->comm->devrState; + const ptrdiff_t stride = (ptrdiff_t)devr->bigSize; + ncclGinAnvilSdmaGPUContext* gpuCtxHost = &ctx->gpuCtxHost[contextId]; + + gpuCtxHost->signals = (uint64_t*)lsaSelf; + + uintptr_t* hostAddrs = (uintptr_t*)calloc((size_t)ctx->nRanks, sizeof(uintptr_t)); + if (!hostAddrs) { + gpuCtxHost->signals = nullptr; return ncclSystemError; } + for (int pe = 0; pe < ctx->nRanks; pe++) { + hostAddrs[pe] = (uintptr_t)lsaSelf + static_cast(pe - ctx->rank) * stride; + } + + uintptr_t selfExpected = (uintptr_t)lsaSelf; + if (hostAddrs[ctx->rank] != selfExpected) { + WARN("GIN anvil-sdma: signal_remote_addrs[self=%d]=%#lx != signals=%#lx (lsaSelf=%#lx stride=%zd)", ctx->rank, + (unsigned long)hostAddrs[ctx->rank], (unsigned long)selfExpected, (unsigned long)lsaSelf, (long)stride); + } - if (ctx->signal_remote_addrs_dev) CUDACHECKIGNORE(hipFree(ctx->signal_remote_addrs_dev)); - if (hipMalloc(&ctx->signal_remote_addrs_dev, sizeof(uintptr_t) * (size_t)ctx->nRanks) != hipSuccess || - hipMemcpy(ctx->signal_remote_addrs_dev, hostAddrs, sizeof(uintptr_t) * (size_t)ctx->nRanks, + if (hipMalloc(&ctx->signal_remote_addrs_dev[contextId], sizeof(uintptr_t) * (size_t)ctx->nRanks) != hipSuccess || + hipMemcpy(ctx->signal_remote_addrs_dev[contextId], hostAddrs, sizeof(uintptr_t) * (size_t)ctx->nRanks, hipMemcpyHostToDevice) != hipSuccess) { free(hostAddrs); + if (ctx->signal_remote_addrs_dev[contextId]) { + CUDACHECKIGNORE(hipFree(ctx->signal_remote_addrs_dev[contextId])); + ctx->signal_remote_addrs_dev[contextId] = nullptr; + } + gpuCtxHost->signals = nullptr; return ncclSystemError; } - ctx->gpuCtxHost.signal_remote_addrs = ctx->signal_remote_addrs_dev; + gpuCtxHost->signal_remote_addrs = ctx->signal_remote_addrs_dev[contextId]; if (ginAnvilSignalDebugEnabled()) { for (int pe = 0; pe < ctx->nRanks; pe++) { @@ -783,64 +856,132 @@ static ncclResult_t ginAnvilRegisterLsaSignals(ginAnvilGinCtx* ctx, void* lsaSel uintptr_t remote0 = hostAddrs[0]; uintptr_t remoteSelf = hostAddrs[ctx->rank]; free(hostAddrs); - ctx->signalsBound = true; - if (hipMemcpy(ctx->gpuCtxDev, &ctx->gpuCtxHost, sizeof(ncclGinAnvilSdmaGPUContext), hipMemcpyHostToDevice) != + if (hipMemcpy(ctx->gpuCtxDev + contextId, gpuCtxHost, sizeof(ncclGinAnvilSdmaGPUContext), hipMemcpyHostToDevice) != hipSuccess) { + CUDACHECKIGNORE(hipFree(ctx->signal_remote_addrs_dev[contextId])); + ctx->signal_remote_addrs_dev[contextId] = nullptr; + gpuCtxHost->signal_remote_addrs = nullptr; + gpuCtxHost->signals = nullptr; return ncclSystemError; } INFO(NCCL_INIT, - "GIN anvil-sdma: bound LSA signals slot=%d signals=%p bytes=%zu rank=%d lsaSelf=%d lsaSize=%d " + "GIN anvil-sdma: bound LSA signals context=%d slot=%d signals=%p bytes=%zu rank=%d lsaSelf=%d lsaSize=%d " "stride=%zu remote[0]=%#lx remote[self]=%#lx", - ctx->signalSlot, lsaSelf, bytes, ctx->rank, devr->lsaSelf, devr->lsaSize, (size_t)devr->bigSize, + contextId, signalSlot, lsaSelf, bytes, ctx->rank, devr->lsaSelf, devr->lsaSize, (size_t)devr->bigSize, (unsigned long)remote0, (unsigned long)remoteSelf); - // [GIN-CONN-CHECK] Validate peer signal connectivity once per comm (on the first - // signal bind that is eligible for the gate). Detects the intermittent gfx950 - // cuMem-VMM peer-map fault and fails loudly here instead of letting the first - // collective hang forever. - bool alreadyChecked = false; - { - std::lock_guard lock(pluginMutex); - alreadyChecked = ginAnvilConnCheckedComms.count(comm) != 0; - } - if (!alreadyChecked) { - NCCLCHECK(ginAnvilCheckSignalConnectivity(ctx, lsaSelf)); - } - return ncclSuccess; } ncclResult_t ncclGinAnvilBindResourceWindowSignals(struct ncclComm* comm, void* resourceUserPtr, size_t arenaByteOffset, - int nContexts, int nSignalsPerContext) { - if (!comm || !resourceUserPtr || nContexts < 1 || nSignalsPerContext < 1) return ncclInvalidArgument; + int nSignalSlots, int nSignalsPerContext) { + if (!comm || !resourceUserPtr || nSignalSlots < 1 || nSignalsPerContext < 1) return ncclInvalidArgument; + + void* spanLocalPtr = (char*)resourceUserPtr + arenaByteOffset; + void* spanLsaSelf = nullptr; + NCCLCHECK(ncclDevrGetLsaSelfAddr(&comm->devrState, spanLocalPtr, &spanLsaSelf)); + if (spanLsaSelf == nullptr) { + WARN("GIN anvil-sdma: could not resolve LSA flat addr for resource-window signal span at %p", spanLocalPtr); + ginAnvilPendingClear(comm); + return ncclSystemError; + } + + // Snapshot the pending list under pluginMutex. ginAnvilPendingAdd / Remove / + // Clear all take that lock and it is not recursive, so the walks below cannot + // hold it across ginAnvilRegisterSignalSpan or ginAnvilPendingClear. find() + // also avoids default-inserting an empty list for a comm with nothing pending. + std::vector pending; + { + std::lock_guard lock(pluginMutex); + auto it = g_pendingByComm.find(comm); + if (it != g_pendingByComm.end()) { + for (GinAnvilPendingEntry* e = it->second; e != nullptr; e = e->next) { + if (e->ctx) pending.push_back(e->ctx); + } + } + } + + const size_t spanBytes = (size_t)nSignalSlots * (size_t)nSignalsPerContext * sizeof(uint64_t); + int spanRanks = 0; + int spanGinRank = 0; + for (ginAnvilGinCtx* ctx : pending) { + if (ctx->nSignals > 0) { + spanRanks = ctx->nRanks; + spanGinRank = ctx->rank; + break; + } + } + if (spanRanks <= 0) { + // Non-anvil GIN backends still call this bind from ncclDevCommCreate whenever + // signals and a resource window exist. An empty pending list means nothing + // anvil-owned needs binding; succeed so type-5 (and similar) bring-up keeps working. + ginAnvilPendingClear(comm); + return ncclSuccess; + } ncclResult_t ret = ncclSuccess; int slot = 0; - for (GinAnvilPendingEntry* e = g_pendingByComm[comm]; e != nullptr; e = e->next) { - ginAnvilGinCtx* ctx = e->ctx; + ginAnvilGinCtx* connCheckCtx = nullptr; + void* connCheckLsaSelf = nullptr; + NCCLCHECKGOTO(ginAnvilRegisterSignalSpan(comm, spanLsaSelf, spanBytes, spanRanks, spanGinRank), ret, fail); + for (ginAnvilGinCtx* ctx : pending) { if (ctx->nSignals <= 0) continue; - ctx->signalSlot = slot++; - if (ctx->signalSlot >= nContexts) { - WARN("GIN anvil-sdma: signal slot %d out of range (nContexts=%d)", ctx->signalSlot, nContexts); - ginAnvilPendingClear(comm); - return ncclInvalidArgument; - } + ctx->signalSpanLsaSelf = spanLsaSelf; + ginAnvilSignalSpanRefInc(spanLsaSelf); + for (int contextId = 0; contextId < ctx->nContexts; contextId++, slot++) { + if (slot >= nSignalSlots) { + WARN("GIN anvil-sdma: signal slot %d out of range (nSignalSlots=%d)", slot, nSignalSlots); + ret = ncclInvalidArgument; + goto fail; + } + + size_t off = arenaByteOffset + (size_t)slot * (size_t)nSignalsPerContext * sizeof(uint64_t); + void* localPtr = (char*)resourceUserPtr + off; + void* lsaSelf = nullptr; + NCCLCHECKGOTO(ncclDevrGetLsaSelfAddr(&comm->devrState, localPtr, &lsaSelf), ret, fail); + if (lsaSelf == nullptr) { + WARN("GIN anvil-sdma: could not resolve LSA flat addr for resource-window signals at %p", localPtr); + ret = ncclSystemError; + goto fail; + } - size_t off = arenaByteOffset + (size_t)ctx->signalSlot * (size_t)nSignalsPerContext * sizeof(uint64_t); - void* localPtr = (char*)resourceUserPtr + off; - void* lsaSelf = nullptr; - NCCLCHECK(ncclDevrGetLsaSelfAddr(&comm->devrState, localPtr, &lsaSelf)); - if (lsaSelf == nullptr) { - WARN("GIN anvil-sdma: could not resolve LSA flat addr for resource-window signals at %p", localPtr); - ginAnvilPendingClear(comm); - return ncclSystemError; + size_t bytes = (size_t)ctx->nSignals * sizeof(uint64_t); + NCCLCHECKGOTO(ginAnvilBindContextSignals(ctx, contextId, slot, lsaSelf, bytes), ret, fail); + if (connCheckCtx == nullptr) { + connCheckCtx = ctx; + connCheckLsaSelf = lsaSelf; + } } + } - size_t bytes = (size_t)ctx->nSignals * sizeof(uint64_t); - NCCLCHECKGOTO(ginAnvilRegisterLsaSignals(ctx, lsaSelf, bytes), ret, fail); + // [GIN-CONN-CHECK] Validate peer signal connectivity once per comm. Detects the + // intermittent gfx950 cuMem-VMM peer-map fault and fails loudly here instead of + // letting the first collective hang forever. This runs per span, not per logical + // context: its setup allgather is collective over the LSA team, and the bypass + // and skip paths return before the comm is marked, so a per-context call would + // repeat that allgather nContexts times. + if (connCheckCtx != nullptr) { + bool alreadyChecked = false; + { + std::lock_guard lock(pluginMutex); + alreadyChecked = ginAnvilConnCheckedComms.count(comm) != 0; + } + if (!alreadyChecked) { + NCCLCHECKGOTO( + ginAnvilCheckSignalConnectivity(connCheckCtx, connCheckLsaSelf, connCheckCtx->signal_remote_addrs_dev[0]), + ret, fail); + } } fail: + if (ret != ncclSuccess) { + for (ginAnvilGinCtx* ctx : pending) { + if (ctx->signalSpanLsaSelf == spanLsaSelf) { + ginAnvilSignalSpanRefDec(spanLsaSelf); + ctx->signalSpanLsaSelf = nullptr; + } + } + } ginAnvilPendingClear(comm); return ret; } @@ -850,14 +991,15 @@ static ncclResult_t ginAnvilCreateContext(void* collComm, ncclGinConfig_t* confi ginAnvilCollCtx* cctx = (ginAnvilCollCtx*)collComm; ncclResult_t ret = ncclSuccess; auto* ctx = new ginAnvilGinCtx{}; + size_t threshold = 0; + ctx->signalSpanLsaSelf = nullptr; ctx->nRanks = cctx->nranks; ctx->rank = cctx->rank; + ctx->nContexts = config->nContexts > 0 ? config->nContexts : 1; ctx->nSignals = config->nSignals; ctx->nCounters = config->nCounters; ctx->comm = cctx->comm; - ctx->signalSlot = -1; // assigned during ncclGinAnvilBindResourceWindowSignals ctx->hasError = false; - ctx->signalsBound = false; ctx->gpu_queue_handles = cctx->gpu_queue_handles; ctx->sdma_dirty_d = cctx->sdma_dirty_d; ctx->numChannels = cctx->numChannels; @@ -875,102 +1017,89 @@ static ncclResult_t ginAnvilCreateContext(void* collComm, ncclGinConfig_t* confi ctx->devHandle->netDeviceVersion = NCCL_GIN_ANVIL_SDMA_NET_VERSION; ctx->devHandle->needsProxyProgress = 0; - if (hipMalloc(&ctx->gpuCtxDev, sizeof(ncclGinAnvilSdmaGPUContext)) != hipSuccess) { + ctx->gpuCtxHost = new ncclGinAnvilSdmaGPUContext[(size_t)ctx->nContexts]{}; + ctx->signal_remote_addrs_dev = (uintptr_t**)calloc((size_t)ctx->nContexts, sizeof(uintptr_t*)); + if (!ctx->signal_remote_addrs_dev) { + ret = ncclSystemError; + goto fail; + } + if (hipMalloc(&ctx->gpuCtxDev, (size_t)ctx->nContexts * sizeof(ncclGinAnvilSdmaGPUContext)) != hipSuccess) { ret = ncclSystemError; goto fail; } - - memset(&ctx->gpuCtxHost, 0, sizeof(ncclGinAnvilSdmaGPUContext)); - ctx->gpuCtxHost.layoutMagic = NCCL_GIN_ANVIL_SDMA_LAYOUT_MAGIC; - ctx->gpuCtxHost.nRanks = ctx->nRanks; - ctx->gpuCtxHost.rank = ctx->rank; - ctx->gpuCtxHost.nSignals = config->nSignals; - ctx->gpuCtxHost.nCounters = config->nCounters; - ctx->gpuCtxHost.numChannels = ctx->numChannels; - ctx->gpuCtxHost.sdmaChannel = 0; - ctx->gpuCtxHost.sdmaChannelStride = ctx->sdmaChannelStride; - ctx->gpuCtxHost.queueHandles = ctx->gpu_queue_handles; - ctx->gpuCtxHost.sdmaDirty = ctx->sdma_dirty_d; - { - const size_t thr = ginAnvilSdmaThresholdFromEnv(); - ctx->gpuCtxHost.sdmaThreshold = - (thr > (size_t)std::numeric_limits::max()) ? std::numeric_limits::max() - : (uint32_t)thr; - } - ctx->gpuCtxHost.fusedSdmaSignal = ginAnvilFusedSignalFromEnv(); - ctx->gpuCtxHost.ipcAgentFence = ginAnvilIpcAgentFenceFromEnv(); - ctx->gpuCtxHost.ipcSignalPeer = ginAnvilIpcSignalPeerFromEnv(); - ctx->gpuCtxHost.signals = nullptr; - ctx->gpuCtxHost.signal_remote_addrs = nullptr; - ctx->signal_remote_addrs_dev = nullptr; if (config->nCounters > 0) { - if (hipExtMallocWithFlags((void**)&ctx->gpuCtxHost.counters, sizeof(uint64_t) * config->nCounters, - hipDeviceMallocFinegrained) != hipSuccess) { + const size_t countersBytes = sizeof(uint64_t) * (size_t)config->nCounters * (size_t)ctx->nContexts; + if (hipExtMallocWithFlags((void**)&ctx->counters, countersBytes, hipDeviceMallocFinegrained) != hipSuccess) { ret = ncclSystemError; goto fail; } - if (hipMemset(ctx->gpuCtxHost.counters, 0, sizeof(uint64_t) * config->nCounters) != hipSuccess) { + if (hipMemset(ctx->counters, 0, countersBytes) != hipSuccess) { ret = ncclSystemError; goto fail; } } - ncclGinAnvilIpcTableGetDevice(&ctx->gpuCtxHost.ipcTable, &ctx->gpuCtxHost.ipcTableCount); + threshold = ginAnvilSdmaThresholdFromEnv(); + for (int contextId = 0; contextId < ctx->nContexts; contextId++) { + ncclGinAnvilSdmaGPUContext* gpuCtx = &ctx->gpuCtxHost[contextId]; + gpuCtx->layoutMagic = NCCL_GIN_ANVIL_SDMA_LAYOUT_MAGIC; + gpuCtx->nRanks = ctx->nRanks; + gpuCtx->rank = ctx->rank; + gpuCtx->nSignals = config->nSignals; + gpuCtx->nCounters = config->nCounters; + gpuCtx->numChannels = ctx->numChannels; + gpuCtx->sdmaChannel = 0; + gpuCtx->sdmaChannelStride = ctx->sdmaChannelStride; + gpuCtx->queueHandles = ctx->gpu_queue_handles; + gpuCtx->sdmaDirty = ctx->sdma_dirty_d; + gpuCtx->sdmaThreshold = (threshold > (size_t)std::numeric_limits::max()) + ? std::numeric_limits::max() + : (uint32_t)threshold; + gpuCtx->fusedSdmaSignal = ginAnvilFusedSignalFromEnv(); + gpuCtx->ipcAgentFence = ginAnvilIpcAgentFenceFromEnv(); + gpuCtx->ipcSignalPeer = ginAnvilIpcSignalPeerFromEnv(); + gpuCtx->counters = + ctx->counters != nullptr ? ctx->counters + (size_t)contextId * (size_t)config->nCounters : nullptr; + ncclGinAnvilIpcTableGetDevice(&gpuCtx->ipcTable, &gpuCtx->ipcTableCount); + } - if (hipMemcpy(ctx->gpuCtxDev, &ctx->gpuCtxHost, sizeof(ncclGinAnvilSdmaGPUContext), hipMemcpyHostToDevice) != - hipSuccess) { + if (hipMemcpy(ctx->gpuCtxDev, ctx->gpuCtxHost, + (size_t)ctx->nContexts * sizeof(ncclGinAnvilSdmaGPUContext), hipMemcpyHostToDevice) != hipSuccess) { WARN("GIN anvil-sdma: hipMemcpy gpu context failed"); ret = ncclSystemError; goto fail; } - ncclGinAnvilIpcTableTrackContext(&ctx->gpuCtxHost, ctx->gpuCtxDev); + for (int contextId = 0; contextId < ctx->nContexts; contextId++) { + ncclGinAnvilIpcTableTrackContext(&ctx->gpuCtxHost[contextId], &ctx->gpuCtxDev[contextId]); + } if (config->nSignals > 0) { ginAnvilPendingAdd(cctx->comm, ctx); } ctx->devHandle->handle = ctx->gpuCtxDev; - ctx->devHandle->size = sizeof(ncclGinAnvilSdmaGPUContext); + ctx->devHandle->size = (size_t)ctx->nContexts * sizeof(ncclGinAnvilSdmaGPUContext); *outGinCtx = ctx; *outDevHandle = ctx->devHandle; INFO(NCCL_INIT, - "GIN anvil-sdma: context created (v%d, %d signals, %d counters, signalSlot=%d, sdmaThreshold=%u, " + "GIN anvil-sdma: context created (v%d, %d logical contexts, %d signals, %d counters, sdmaThreshold=%u, " "spread=%d, fusedSignal=%u)", - NCCL_GIN_ANVIL_SDMA_NET_VERSION, config->nSignals, config->nCounters, ctx->signalSlot, - ctx->gpuCtxHost.sdmaThreshold, ctx->gpuCtxHost.sdmaChannelStride, ctx->gpuCtxHost.fusedSdmaSignal); + NCCL_GIN_ANVIL_SDMA_NET_VERSION, ctx->nContexts, config->nSignals, config->nCounters, + ctx->gpuCtxHost[0].sdmaThreshold, ctx->gpuCtxHost[0].sdmaChannelStride, ctx->gpuCtxHost[0].fusedSdmaSignal); return ncclSuccess; fail: - if (ctx) { - if (ctx->comm) ginAnvilPendingRemove(ctx->comm, ctx); - if (ctx->signalsBound && ctx->gpuCtxHost.signals) { - (void)ncclGinAnvilIpcTableUnregister(ctx->gpuCtxHost.signals); - } - if (ctx->signal_remote_addrs_dev) CUDACHECKIGNORE(hipFree(ctx->signal_remote_addrs_dev)); - if (ctx->gpuCtxHost.counters) CUDACHECKIGNORE(hipFree(ctx->gpuCtxHost.counters)); - if (ctx->gpuCtxDev) CUDACHECKIGNORE(hipFree(ctx->gpuCtxDev)); - free(ctx->devHandle); - delete ctx; - } + if (ctx) ginAnvilFreeGinCtx(ctx); return ret; } static ncclResult_t ginAnvilDestroyContext(void* ginCtx) { ginAnvilGinCtx* ctx = (ginAnvilGinCtx*)ginCtx; if (!ctx) return ncclSuccess; - if (ctx->comm) ginAnvilPendingRemove(ctx->comm, ctx); - ncclGinAnvilIpcTableUntrackContext(&ctx->gpuCtxHost); - if (ctx->signalsBound && ctx->gpuCtxHost.signals) { - (void)ncclGinAnvilIpcTableUnregister(ctx->gpuCtxHost.signals); - } - if (ctx->signal_remote_addrs_dev) CUDACHECKIGNORE(hipFree(ctx->signal_remote_addrs_dev)); - if (ctx->gpuCtxHost.counters) CUDACHECKIGNORE(hipFree(ctx->gpuCtxHost.counters)); - if (ctx->gpuCtxDev) CUDACHECKIGNORE(hipFree(ctx->gpuCtxDev)); - free(ctx->devHandle); - delete ctx; + ginAnvilFreeGinCtx(ctx); return ncclSuccess; } diff --git a/projects/rccl/src/include/gin/gin_host_anvil_sdma.h b/projects/rccl/src/include/gin/gin_host_anvil_sdma.h index d18f4a2dbb9b..c119af00e5db 100644 --- a/projects/rccl/src/include/gin/gin_host_anvil_sdma.h +++ b/projects/rccl/src/include/gin/gin_host_anvil_sdma.h @@ -15,7 +15,7 @@ extern ncclGin_t ncclGinAnvilSdmaPlugin; struct ncclComm; ncclResult_t ncclGinAnvilBindResourceWindowSignals(struct ncclComm* comm, void* resourceUserPtr, size_t arenaByteOffset, - int nContexts, int nSignalsPerContext); + int nSignalSlots, int nSignalsPerContext); void ncclGinAnvilSetInitContext(void* initCtx, struct ncclComm* comm); diff --git a/projects/rccl/src/include/nccl_device/gin/anvil_sdma/gin_anvil_sdma.h b/projects/rccl/src/include/nccl_device/gin/anvil_sdma/gin_anvil_sdma.h index 8f0efee0f209..40cafc3201b8 100644 --- a/projects/rccl/src/include/nccl_device/gin/anvil_sdma/gin_anvil_sdma.h +++ b/projects/rccl/src/include/nccl_device/gin/anvil_sdma/gin_anvil_sdma.h @@ -35,6 +35,12 @@ NCCL_DEVICE_INLINE bool anvilCtxValid(ncclGinAnvilSdmaGPUContext* rsCtx) { return rsCtx != nullptr && loadConst(&rsCtx->layoutMagic) == NCCL_GIN_ANVIL_SDMA_LAYOUT_MAGIC; } +// Logical contexts index distinct GPU-context slots. Counters are sliced per +// context (countersDev + contextId * nCounters), matching proxy/gdaki. +NCCL_DEVICE_INLINE ncclGinAnvilSdmaGPUContext* anvilGpuCtx(ncclGinCtx ctx) { + return &reinterpret_cast(ctx.handle)[ctx.contextId]; +} + NCCL_DEVICE_INLINE void* resolveRemotePeerVa(ncclGinAnvilSdmaGPUContext* rsCtx, ncclGinAnvilSdmaMemHandle* mh, int peer, size_t off) { ptrdiff_t stride = loadConst(&mh->vmmStride); @@ -183,7 +189,7 @@ struct ncclGinApi_Put { if (coop.thread_rank() != 0) return; - ncclGinAnvilSdmaGPUContext* rsCtx = (ncclGinAnvilSdmaGPUContext*)ctx.handle; + ncclGinAnvilSdmaGPUContext* rsCtx = nccl::gin::anvil::detail::anvilGpuCtx(ctx); if (!anvilCtxValid(rsCtx)) return; const int blockId = blockIdx.x + blockIdx.y * gridDim.x; @@ -298,7 +304,7 @@ struct ncclGinApi_PutValue { if (coop.thread_rank() != 0) return; - ncclGinAnvilSdmaGPUContext* rsCtx = (ncclGinAnvilSdmaGPUContext*)ctx.handle; + ncclGinAnvilSdmaGPUContext* rsCtx = nccl::gin::anvil::detail::anvilGpuCtx(ctx); if (!anvilCtxValid(rsCtx)) return; const int blockId = blockIdx.x + blockIdx.y * gridDim.x; ncclGinAnvilSdmaMemHandle* dstMh = (ncclGinAnvilSdmaMemHandle*)dstWin; @@ -367,7 +373,7 @@ struct ncclGinApi_PutValue { template <> struct ncclGinApi_GetCounterPtr { NCCL_DEVICE_INLINE static ncclGinOffsetPtr call(ncclGinCtx ctx, ncclGinCounter_t counterId) { - ncclGinAnvilSdmaGPUContext* rsCtx = (ncclGinAnvilSdmaGPUContext*)ctx.handle; + ncclGinAnvilSdmaGPUContext* rsCtx = nccl::gin::anvil::detail::anvilGpuCtx(ctx); if (!nccl::gin::anvil::detail::anvilCtxValid(rsCtx)) return {nullptr, 0}; return {nccl::utility::loadConst(&rsCtx->counters) + counterId, 0}; } @@ -376,7 +382,7 @@ struct ncclGinApi_GetCounterPtr { template <> struct ncclGinApi_ResetCounter { NCCL_DEVICE_INLINE static void call(ncclGinCtx ctx, ncclGinCounter_t counterId) { - ncclGinAnvilSdmaGPUContext* rsCtx = (ncclGinAnvilSdmaGPUContext*)ctx.handle; + ncclGinAnvilSdmaGPUContext* rsCtx = nccl::gin::anvil::detail::anvilGpuCtx(ctx); if (!nccl::gin::anvil::detail::anvilCtxValid(rsCtx)) return; uint64_t* counters = nccl::utility::loadConst(&rsCtx->counters); if (counters != nullptr) counters[counterId] = 0; @@ -388,7 +394,7 @@ struct ncclGinApi_GetSignalPtr { NCCL_DEVICE_INLINE static ncclGinOffsetPtr call(ncclGinCtx ctx, ncclGinSignal_t signalId) { using nccl::gin::anvil::detail::anvilCtxValid; using nccl::utility::loadConst; - ncclGinAnvilSdmaGPUContext* rsCtx = (ncclGinAnvilSdmaGPUContext*)ctx.handle; + ncclGinAnvilSdmaGPUContext* rsCtx = nccl::gin::anvil::detail::anvilGpuCtx(ctx); assert(anvilCtxValid(rsCtx)); uint64_t* signals = loadConst(&rsCtx->signals); assert(signals != nullptr); @@ -399,7 +405,7 @@ struct ncclGinApi_GetSignalPtr { template <> struct ncclGinApi_ResetSignal { NCCL_DEVICE_INLINE static void call(ncclGinCtx ctx, ncclGinSignalDescriptor signal) { - ncclGinAnvilSdmaGPUContext* rsCtx = (ncclGinAnvilSdmaGPUContext*)ctx.handle; + ncclGinAnvilSdmaGPUContext* rsCtx = nccl::gin::anvil::detail::anvilGpuCtx(ctx); if (!nccl::gin::anvil::detail::anvilCtxValid(rsCtx)) return; if (signal.type == NCCL_GIN_SIGNAL_TYPE_INDEXED) { uint64_t* signals = nccl::utility::loadConst(&rsCtx->signals); @@ -418,7 +424,7 @@ struct ncclGinApi_Flush { (void)ord; (void)abortFlag; using nccl::utility::loadConst; - ncclGinAnvilSdmaGPUContext* rsCtx = (ncclGinAnvilSdmaGPUContext*)ctx.handle; + ncclGinAnvilSdmaGPUContext* rsCtx = nccl::gin::anvil::detail::anvilGpuCtx(ctx); if (!nccl::gin::anvil::detail::anvilCtxValid(rsCtx)) { __threadfence_system(); return; @@ -475,7 +481,7 @@ struct ncclGinApi_Get { if (coop.thread_rank() != 0 || bytes == 0) return; - ncclGinAnvilSdmaGPUContext* rsCtx = (ncclGinAnvilSdmaGPUContext*)ctx.handle; + ncclGinAnvilSdmaGPUContext* rsCtx = nccl::gin::anvil::detail::anvilGpuCtx(ctx); if (!anvilCtxValid(rsCtx)) return; const int blockId = blockIdx.x + blockIdx.y * gridDim.x; ncclGinAnvilSdmaMemHandle* remoteMh = (ncclGinAnvilSdmaMemHandle*)remoteWin; @@ -519,7 +525,7 @@ struct ncclGinApi_FlushAsync { auto* request = reinterpret_cast(outRequest); request->complete = 0; - ncclGinAnvilSdmaGPUContext* rsCtx = (ncclGinAnvilSdmaGPUContext*)ctx.handle; + ncclGinAnvilSdmaGPUContext* rsCtx = nccl::gin::anvil::detail::anvilGpuCtx(ctx); if (nccl::gin::anvil::detail::anvilCtxValid(rsCtx)) { uint64_t* dirtyPtr = loadConst(&rsCtx->sdmaDirty); uint64_t dirty = diff --git a/projects/rccl/test/device/GinAnvilSdmaTemplate_test.cpp b/projects/rccl/test/device/GinAnvilSdmaTemplate_test.cpp index 2289863f6f0f..c2f266d2e347 100644 --- a/projects/rccl/test/device/GinAnvilSdmaTemplate_test.cpp +++ b/projects/rccl/test/device/GinAnvilSdmaTemplate_test.cpp @@ -396,6 +396,118 @@ TEST_F(GinAnvilSdmaTemplateTest, CounterSignal_GetReset) { EXPECT_EQ(d_outSig.download(), 0ULL); } +// AICOMRCCL-2339: the Anvil device dispatch must select the GPU context +// indexed by ncclGinCtx::contextId instead of always using array element zero. +__global__ void kernelSignalContextSelection(ncclGinAnvilSdmaGPUContext* contexts) { + if (threadIdx.x != 0) return; + ncclGinCtx ginCtx{}; + ginCtx.handle = contexts; + + ginCtx.contextId = 0; + ncclGinApi_GetSignalPtr::call(ginCtx, 0).ptr[0] = 11; + ncclGinApi_GetCounterPtr::call(ginCtx, 0).ptr[0] = 101; + + ginCtx.contextId = 1; + ncclGinApi_GetSignalPtr::call(ginCtx, 0).ptr[0] = 22; + ncclGinApi_ResetSignal::call( + ginCtx, ncclGinSignalDescriptor{NCCL_GIN_SIGNAL_TYPE_INDEXED, {.indexedSignal = {.signalId = 0}}}); + ncclGinApi_ResetCounter::call(ginCtx, 0); + ncclGinApi_GetCounterPtr::call(ginCtx, 0).ptr[0] = 202; +} + +TEST_F(GinAnvilSdmaTemplateTest, SignalApis_SelectLogicalContext) { + DeviceBuffer d_signals(2); + DeviceBuffer d_counters(2); + d_signals.zero(); + d_counters.zero(); + ncclGinAnvilSdmaGPUContext hostCtx[2]{}; + for (int i = 0; i < 2; i++) { + hostCtx[i].layoutMagic = NCCL_GIN_ANVIL_SDMA_LAYOUT_MAGIC; + hostCtx[i].signals = d_signals.ptr + i; + hostCtx[i].counters = d_counters.ptr + i; + hostCtx[i].nSignals = 1; + hostCtx[i].nCounters = 1; + } + DeviceBuffer d_contexts(2); + d_contexts.copyFrom(hostCtx, 2); + + kernelSignalContextSelection<<<1, 1>>>(d_contexts.ptr); + syncAndCheck(); + auto signals = d_signals.copyTo(); + auto counters = d_counters.copyTo(); + EXPECT_EQ(signals[0], 11ULL) << "ResetSignal must not clear context 0 when invoked on contextId=1"; + EXPECT_EQ(signals[1], 0ULL) << "ResetSignal on contextId=1 must clear context 1 signal cell"; + EXPECT_EQ(counters[0], 101ULL); + EXPECT_EQ(counters[1], 202ULL); +} + +// AICOMRCCL-2339: PutValue must honor ncclGinCtx::contextId when resolving signals. +__global__ void kernelPutValueContextSelection(ncclGinAnvilSdmaGPUContext* contexts, ncclGinAnvilSdmaMemHandle* dstMh) { + if (threadIdx.x != 0) return; + ncclGinCtx ginCtx{}; + ginCtx.handle = contexts; + ginCtx.nRanks = 2; + ginCtx.rank = 0; + ginCtx.contextId = 1; + + ncclGinSignalDescriptor sig{}; + sig.type = NCCL_GIN_SIGNAL_TYPE_INDEXED; + sig.indexedSignal.signalId = 0; + ncclGinApi_PutValue::call( + ginCtx, ncclCoopThread{}, 1, reinterpret_cast(dstMh), 0, static_cast(0x11), sig, + ncclGinSignalInc, 1, false, nullptr, cuda::thread_scope_system, cuda::thread_scope_system); +} + +TEST_F(GinAnvilSdmaTemplateTest, PutValue_SelectLogicalContext) { + DeviceBuffer d_signals(2); + DeviceBuffer d_dst(8); + DeviceBuffer d_remoteAddrs0(2); + DeviceBuffer d_remoteAddrs1(2); + d_signals.zero(); + d_dst.zero(); + + // Distinct peer-1 fallback entries so reverting anvilGpuCtx to handle[0] + // would write context 0's cell and fail the assertions below. + uintptr_t remoteAddrs0[2] = {0, reinterpret_cast(d_signals.ptr + 0)}; + uintptr_t remoteAddrs1[2] = {0, reinterpret_cast(d_signals.ptr + 1)}; + d_remoteAddrs0.copyFrom(remoteAddrs0, 2); + d_remoteAddrs1.copyFrom(remoteAddrs1, 2); + + DeviceBuffer d_entry(1); + ncclGinAnvilIpcBufEntry ipcEntry{}; + ipcEntry.local_base = reinterpret_cast(d_dst.ptr); + ipcEntry.length = 4096; + ipcEntry.remote_bases[1] = reinterpret_cast(d_dst.ptr); + d_entry.upload(ipcEntry); + + ncclGinAnvilSdmaMemHandle dstMh{}; + dstMh.baseAddr = reinterpret_cast(d_dst.ptr); + DeviceBuffer d_dstMh(1); + d_dstMh.upload(dstMh); + + ncclGinAnvilSdmaGPUContext hostCtx[2]{}; + for (int i = 0; i < 2; i++) { + hostCtx[i].layoutMagic = NCCL_GIN_ANVIL_SDMA_LAYOUT_MAGIC; + hostCtx[i].signals = d_signals.ptr + i; + hostCtx[i].signal_remote_addrs = (i == 0) ? d_remoteAddrs0.ptr : d_remoteAddrs1.ptr; + hostCtx[i].ipcTable = d_entry.ptr; + hostCtx[i].ipcTableCount = 1; + hostCtx[i].nRanks = 2; + hostCtx[i].rank = 0; + hostCtx[i].sdmaThreshold = 4096; + hostCtx[i].fusedSdmaSignal = 0; + } + + DeviceBuffer d_contexts(2); + d_contexts.copyFrom(hostCtx, 2); + + kernelPutValueContextSelection<<<1, 1>>>(d_contexts.ptr, d_dstMh.ptr); + syncAndCheck(); + auto signals = d_signals.copyTo(); + EXPECT_EQ(signals[0], 0ULL) << "context 0 signal stripe must stay untouched"; + EXPECT_EQ(signals[1], 1ULL) << "PutValue on contextId=1 must increment context 1 signal cell"; +} + // H10: invalid ctx on getters returns nullptr / no-op. __global__ void kernelInvalidCtxApis(bool* ok) { ncclGinCtx ginCtx{}; diff --git a/projects/rccl/test/gin/GinAnvilPlugin_test.cpp b/projects/rccl/test/gin/GinAnvilPlugin_test.cpp index 817a7173f40b..24fe45202325 100644 --- a/projects/rccl/test/gin/GinAnvilPlugin_test.cpp +++ b/projects/rccl/test/gin/GinAnvilPlugin_test.cpp @@ -364,6 +364,65 @@ TEST_F(GinAnvilPluginTest, BindSignals_Success) { plugin_.finalize(ictx); } +// AICOMRCCL-2339: logical contexts on one Anvil connection share SDMA +// queues, but each context must bind a distinct signal stripe. A second +// pending context on the same comm makes slot and contextId diverge: the +// first context's contextId 0 is not slot 0. +TEST_F(GinAnvilPluginTest, BindSignals_LogicalContextsUseDistinctSignalStripes) { + void* ictx = nullptr; + initCtx(&ictx); + void* coll = nullptr; + connectColl(ictx, &coll); + ncclGinConfig_t cfg{}; + cfg.nContexts = 3; + cfg.nSignals = 2; + cfg.nCounters = 5; + void* ginCtx = nullptr; + ncclNetDeviceHandle_v11_t* devHandle = nullptr; + ASSERT_EQ(plugin_.createContext(coll, &cfg, &ginCtx, &devHandle), ncclSuccess); + + ncclGinConfig_t cfgB{}; + cfgB.nContexts = 1; + cfgB.nSignals = 2; + cfgB.nCounters = 5; + void* ginCtxB = nullptr; + ncclNetDeviceHandle_v11_t* devHandleB = nullptr; + ASSERT_EQ(plugin_.createContext(coll, &cfgB, &ginCtxB, &devHandleB), ncclSuccess); + + char arena[8192] = {}; + // Pending list is newest-first: ginCtxB takes slot 0, ginCtx takes slots 1..3. + ASSERT_EQ(ncclGinAnvilBindResourceWindowSignals(mockComm_.get(), arena, 0, 4, 4), ncclSuccess); + ASSERT_EQ(devHandle->size, 3 * sizeof(ncclGinAnvilSdmaGPUContext)); + ASSERT_EQ(devHandleB->size, sizeof(ncclGinAnvilSdmaGPUContext)); + + ncclGinAnvilSdmaGPUContext hostCtx[3]{}; + ncclGinAnvilSdmaGPUContext hostCtxB{}; + ASSERT_EQ(hipMemcpy(hostCtx, devHandle->handle, sizeof(hostCtx), hipMemcpyDeviceToHost), hipSuccess); + ASSERT_EQ(hipMemcpy(&hostCtxB, devHandleB->handle, sizeof(hostCtxB), hipMemcpyDeviceToHost), hipSuccess); + ASSERT_NE(hostCtx[0].signals, nullptr); + ASSERT_NE(hostCtx[1].signals, nullptr); + ASSERT_NE(hostCtx[2].signals, nullptr); + ASSERT_NE(hostCtxB.signals, nullptr); + EXPECT_EQ(hostCtx[1].signals - hostCtx[0].signals, 4); + EXPECT_EQ(hostCtx[2].signals - hostCtx[1].signals, 4); + EXPECT_EQ(hostCtx[0].signals - hostCtxB.signals, 4); + EXPECT_NE(hostCtx[0].signal_remote_addrs, hostCtx[1].signal_remote_addrs); + + // Queue ownership remains at the connection level by design. + EXPECT_EQ(hostCtx[0].queueHandles, hostCtx[1].queueHandles); + EXPECT_EQ(hostCtx[0].sdmaDirty, hostCtx[1].sdmaDirty); + EXPECT_EQ(hostCtx[0].queueHandles, hostCtxB.queueHandles); + EXPECT_NE(hostCtx[0].counters, nullptr); + EXPECT_NE(hostCtx[1].counters, nullptr); + EXPECT_EQ(hostCtx[1].counters - hostCtx[0].counters, 5); + EXPECT_EQ(hostCtx[2].counters - hostCtx[1].counters, 5); + + plugin_.destroyContext(ginCtx); + plugin_.destroyContext(ginCtxB); + plugin_.closeColl(coll); + plugin_.finalize(ictx); +} + // G15: dereg null, progress, queryLastError. TEST_F(GinAnvilPluginTest, Misc_NoOpPaths) { EXPECT_EQ(plugin_.deregMrSym(nullptr, nullptr), ncclSuccess); @@ -661,6 +720,85 @@ TEST_F(GinAnvilPluginTest, ConnCheck_HealthyConnectivitySucceeds) { plugin_.finalize(ictx); } +// AICOMRCCL-2339: the comm-wide collectives in the bind path are per span, not +// per logical context. Binding a context with nContexts=3 must still cost one +// local-base allgather and one conn-check, or ncclDevCommCreate pays a bootstrap +// round trip for every logical context. +TEST_F(GinAnvilPluginTest, BindSignals_CommWideCollectivesRunOncePerSpan) { + void* rawDevLsa = nullptr; + ASSERT_EQ(hipMalloc(&rawDevLsa, sizeof(uint64_t) * 8), hipSuccess); + HipAllocation devLsa(rawDevLsa); + GinAnvilPluginStubs::SetLsaSelfAddr(devLsa.get()); + GinAnvilPluginStubs::SetBootstrapNranks(2); + mockComm_.get()->devrState.lsaSize = 2; + + void* ictx = nullptr; + initCtx(&ictx); + void* coll = nullptr; + connectColl(ictx, &coll, 2); + ncclGinConfig_t cfg{}; + cfg.nContexts = 3; + cfg.nSignals = 2; + void* ginCtx = nullptr; + ncclNetDeviceHandle_v11_t* devHandle = nullptr; + ASSERT_EQ(plugin_.createContext(coll, &cfg, &ginCtx, &devHandle), ncclSuccess); + + // connectColl runs the factory allgather; only the bind path is under test. + const int allGathersBeforeBind = GinAnvilPluginStubs::GetBootstrapAllGatherCalls(); + const int intraNodeBeforeBind = GinAnvilPluginStubs::GetIntraNodeAllGatherCalls(); + + char arena[8192] = {}; + ASSERT_EQ(ncclGinAnvilBindResourceWindowSignals(mockComm_.get(), arena, 0, 3, 2), ncclSuccess); + + EXPECT_EQ(GinAnvilPluginStubs::GetBootstrapAllGatherCalls() - allGathersBeforeBind, 1); + // One conn-check, which costs two LSA allgathers: the setup state exchange and + // the missing-count exchange of its single attempt. + EXPECT_EQ(GinAnvilPluginStubs::GetIntraNodeAllGatherCalls() - intraNodeBeforeBind, 2); + EXPECT_EQ(GinAnvilPluginStubs::GetConnCheckWriteCalls(), 1); + EXPECT_EQ(GinAnvilPluginStubs::GetConnCheckVerifyCalls(), 1); + + // All three contexts are still bound to their own stripe. + ncclGinAnvilSdmaGPUContext hostCtx[3]{}; + ASSERT_EQ(hipMemcpy(hostCtx, devHandle->handle, sizeof(hostCtx), hipMemcpyDeviceToHost), hipSuccess); + EXPECT_EQ(hostCtx[1].signals - hostCtx[0].signals, 2); + EXPECT_EQ(hostCtx[2].signals - hostCtx[1].signals, 2); + + plugin_.destroyContext(ginCtx); + plugin_.closeColl(coll); + plugin_.finalize(ictx); +} + +// The skip path returns before the comm is marked as checked, so a per-context +// conn-check call would repeat its setup allgather nContexts times. +TEST_F(GinAnvilPluginTest, BindSignals_SkippedConnCheckStillRunsOncePerSpan) { + GinAnvilPluginStubs::SetBootstrapNranks(2); + mockComm_.get()->devrState.lsaSize = 4; // != nRanks, so the gate skips + + void* ictx = nullptr; + initCtx(&ictx); + void* coll = nullptr; + connectColl(ictx, &coll, 2); + ncclGinConfig_t cfg{}; + cfg.nContexts = 3; + cfg.nSignals = 2; + void* ginCtx = nullptr; + ncclNetDeviceHandle_v11_t* devHandle = nullptr; + ASSERT_EQ(plugin_.createContext(coll, &cfg, &ginCtx, &devHandle), ncclSuccess); + + const int intraNodeBeforeBind = GinAnvilPluginStubs::GetIntraNodeAllGatherCalls(); + + char arena[8192] = {}; + ASSERT_EQ(ncclGinAnvilBindResourceWindowSignals(mockComm_.get(), arena, 0, 3, 2), ncclSuccess); + + EXPECT_EQ(GinAnvilPluginStubs::GetIntraNodeAllGatherCalls() - intraNodeBeforeBind, 1); + EXPECT_EQ(GinAnvilPluginStubs::GetConnCheckWriteCalls(), 0); + EXPECT_EQ(GinAnvilPluginStubs::GetConnCheckVerifyCalls(), 0); + + plugin_.destroyContext(ginCtx); + plugin_.closeColl(coll); + plugin_.finalize(ictx); +} + TEST_F(GinAnvilPluginTest, ConnCheck_NonNumericEnvValueKeepsGateEnabled) { ScopedEnv enabled("NCCL_GIN_ANVIL_SDMA_CONN_CHECK", "true"); void* rawDevLsa = nullptr; diff --git a/projects/rccl/test/gin/gin_anvil_plugin_test_stubs.cc b/projects/rccl/test/gin/gin_anvil_plugin_test_stubs.cc index af8544f55be1..b19a286b39cf 100644 --- a/projects/rccl/test/gin/gin_anvil_plugin_test_stubs.cc +++ b/projects/rccl/test/gin/gin_anvil_plugin_test_stubs.cc @@ -41,7 +41,10 @@ struct State { int connCheckWriteCalls = 0; int connCheckVerifyCalls = 0; std::vector connCheckWriteStamps; + int bootstrapAllGatherCalls = 0; + int intraNodeAllGatherCalls = 0; void* lsaSelfAddr = reinterpret_cast(0x70001000ULL); + uintptr_t lsaInputBase = 0; }; struct FakeSdmaOpaque { @@ -65,7 +68,12 @@ void SetBootstrapIntResult(const int* values, int count) { void SetFactoryCreateFail(bool fail) { g.factoryCreateFail = fail; } void SetFactoryNullHandles(bool nullHandles) { g.factoryNullHandles = nullHandles; } void SetLsaAddrFail(bool fail) { g.lsaAddrFail = fail; } -void SetLsaSelfAddr(void* addr) { g.lsaSelfAddr = addr; } +void SetLsaSelfAddr(void* addr) { + g.lsaSelfAddr = addr; + // Drop the latched input base so the next resolve is relative to this address, + // not the first arena that happened to be resolved. + g.lsaInputBase = 0; +} void SetConnCheckMissingCalls(int calls) { g.connCheckMissingCalls = calls; } int GetConnCheckWriteCalls() { return g.connCheckWriteCalls; } int GetConnCheckVerifyCalls() { return g.connCheckVerifyCalls; } @@ -77,6 +85,8 @@ unsigned long long GetConnCheckWriteStamp(int call) { const std::vector& GetLastIntraNodeAllGatherRanks() { return g.lastIntraNodeAllGather.ranks; } int GetLastIntraNodeAllGatherRank() { return g.lastIntraNodeAllGather.rank; } int GetLastIntraNodeAllGatherNranks() { return g.lastIntraNodeAllGather.nranks; } +int GetBootstrapAllGatherCalls() { return g.bootstrapAllGatherCalls; } +int GetIntraNodeAllGatherCalls() { return g.intraNodeAllGatherCalls; } const std::vector& GetLastIntraNodeBarrierRanks() { return g.lastIntraNodeBarrier.ranks; } int GetLastIntraNodeBarrierRank() { return g.lastIntraNodeBarrier.rank; } int GetLastIntraNodeBarrierTag() { return g.lastIntraNodeBarrier.tag; } @@ -127,6 +137,7 @@ static ncclResult_t stubIntAllGather(void* allData, int nranks, int size) { ncclResult_t bootstrapAllGather(void* commState, void* allData, int size) { (void)commState; + GinAnvilPluginStubs::g.bootstrapAllGatherCalls++; return stubIntAllGather(allData, GinAnvilPluginStubs::g.bootstrapNranks, size); } @@ -145,6 +156,7 @@ ncclResult_t bootstrapIntraNodeAllGather(void* commState, int* ranks, int rank, GinAnvilPluginStubs::g.lastIntraNodeAllGather.ranks.assign(ranks, ranks + nranks); GinAnvilPluginStubs::g.lastIntraNodeAllGather.rank = rank; GinAnvilPluginStubs::g.lastIntraNodeAllGather.nranks = nranks; + GinAnvilPluginStubs::g.intraNodeAllGatherCalls++; return stubIntAllGather(allData, nranks, size); } @@ -158,12 +170,14 @@ ncclResult_t bootstrapIntraNodeBarrier(void* commState, int* ranks, int rank, in ncclResult_t ncclDevrGetLsaSelfAddr(struct ncclDevrState* devr, void* addr, void** outAddr) { (void)devr; - (void)addr; if (GinAnvilPluginStubs::g.lsaAddrFail) { *outAddr = nullptr; return ncclSuccess; } - *outAddr = GinAnvilPluginStubs::g.lsaSelfAddr; + uintptr_t input = reinterpret_cast(addr); + if (GinAnvilPluginStubs::g.lsaInputBase == 0) GinAnvilPluginStubs::g.lsaInputBase = input; + uintptr_t offset = input - GinAnvilPluginStubs::g.lsaInputBase; + *outAddr = reinterpret_cast(reinterpret_cast(GinAnvilPluginStubs::g.lsaSelfAddr) + offset); return ncclSuccess; } diff --git a/projects/rccl/test/gin/gin_anvil_plugin_test_stubs.h b/projects/rccl/test/gin/gin_anvil_plugin_test_stubs.h index cec0481fa30d..9197b3f70db2 100644 --- a/projects/rccl/test/gin/gin_anvil_plugin_test_stubs.h +++ b/projects/rccl/test/gin/gin_anvil_plugin_test_stubs.h @@ -30,6 +30,8 @@ unsigned long long GetConnCheckWriteStamp(int call); const std::vector& GetLastIntraNodeAllGatherRanks(); int GetLastIntraNodeAllGatherRank(); int GetLastIntraNodeAllGatherNranks(); +int GetBootstrapAllGatherCalls(); +int GetIntraNodeAllGatherCalls(); const std::vector& GetLastIntraNodeBarrierRanks(); int GetLastIntraNodeBarrierRank(); int GetLastIntraNodeBarrierTag(); diff --git a/projects/rccl/test/transport/GinDeviceMPITests.cpp b/projects/rccl/test/transport/GinDeviceMPITests.cpp index cc5c72d9e027..b89248ade5da 100644 --- a/projects/rccl/test/transport/GinDeviceMPITests.cpp +++ b/projects/rccl/test/transport/GinDeviceMPITests.cpp @@ -1973,6 +1973,108 @@ TEST_F(GinMPIDeviceTests, Barrier_TwoRanks) { MPI_Barrier(MPI_COMM_WORLD); } +// AICOMRCCL-2339 regression: each logical context must have independent +// signal cells. Rank 1 deliberately arrives late at the second barrier. Before +// the fix, both contexts incremented one shared cell during the first barrier, +// so rank 0's second wait matched that stale count and returned immediately. +struct AllContextsBarrierObservation { + uint64_t secondBarrierCycles; + uint64_t ctxPeerSignals[2]; + uint64_t* ctxSignalPtrs[2]; +}; + +__global__ void allContextsConsecutiveBarrierKernel( + int rank, uint64_t delayCycles, AllContextsBarrierObservation* observation, struct ncclDevComm devComm) { + ncclGinBarrierSession bar{ + ncclCoopCta(), ncclGinAllContexts(devComm), ncclTeamTagWorld{}, /*barrierIndex=*/0}; + bar.sync(ncclCoopCta(), cuda::memory_order_relaxed, ncclGinFenceLevel::Relaxed); + + if (rank == 1) { + uint64_t start = clock64(); + while (clock64() - start < delayCycles) {} + } + ncclCoopCta().sync(); + + uint64_t start = clock64(); + bar.sync(ncclCoopCta(), cuda::memory_order_relaxed, ncclGinFenceLevel::Relaxed); + if (threadIdx.x != 0) return; + + observation->secondBarrierCycles = clock64() - start; + const int peer = 1 - rank; + const uint32_t sigIndex = devComm.worldGinBarrier.signal0 + peer; + for (int ctx = 0; ctx < 2; ctx++) { + ncclGin gin{devComm, ctx}; + observation->ctxPeerSignals[ctx] = gin.readSignal(sigIndex, 64, cuda::memory_order_relaxed); +#if NCCL_GIN_ANVIL_SDMA_ENABLE + // Derive handle/contextId the same way ncclGinInitCommon does so + // NCCL_GIN_NCONNECTIONS!=1 does not read past ginHandles[0]. + ncclGinCtx ginCtx{}; + ginCtx.handle = gin._ginHandle; + ginCtx.contextId = gin.contextId; + ginCtx.backend = NCCL_NET_DEVICE_GIN_ANVIL_SDMA; + ginCtx.rank = devComm.rank; + ginCtx.nRanks = devComm.nRanks; + observation->ctxSignalPtrs[ctx] = + ncclGinApi_GetSignalPtr::call(ginCtx, sigIndex).ptr; +#endif + } +} + +TEST_F(GinMPIDeviceTests, Barrier_AllContextsConsecutiveSignalsDoNotAlias_SingleNode) { + if (auto reason = ginProxyTestSkipReason(); !reason.empty()) + GTEST_SKIP() << reason; + if (requestedGinType() != NCCL_NET_DEVICE_GIN_ANVIL_SDMA) + GTEST_SKIP() << "AICOMRCCL-2339 is specific to the Anvil-SDMA backend"; + if (!validateTestPrerequisites(/*min_processes=*/2, /*max_processes=*/2)) + GTEST_SKIP() << "Requires exactly 2 ranks"; + if (nodeLocalRanks() != 2) + GTEST_SKIP() << "Requires both ranks on one node"; + + ASSERT_EQ(ncclSuccess, createTestCommunicator()); + ncclComm_t comm = getActiveCommunicator(); + hipStream_t stream = getActiveStream(); + int rank = -1; + ncclCommUserRank(comm, &rank); + + ncclDevCommRequirements reqs = defaultGinReqs(); + reqs.ginContextCount = 2; + reqs.worldGinBarrierCount = 1; + ncclDevComm devComm{}; + ASSERT_MPI_EQ(ncclSuccess, ncclDevCommCreate(comm, &reqs, &devComm)); + auto devCommCleanup = makeScopeGuard([&]() { + (void)ncclDevCommDestroy(comm, &devComm); + }); + ASSERT_GE((int)devComm.ginContextCount, 2); + + AllContextsBarrierObservation* dObservation = nullptr; + ASSERT_MPI_EQ(hipSuccess, hipMalloc(&dObservation, sizeof(AllContextsBarrierObservation))); + auto observationCleanup = makeScopeGuard([&]() { + if (dObservation) (void)hipFree(dObservation); + }); + ASSERT_MPI_EQ(hipSuccess, hipMemset(dObservation, 0, sizeof(AllContextsBarrierObservation))); + + constexpr uint64_t kDelayCycles = 100000000; + MPI_Barrier(MPI_COMM_WORLD); + allContextsConsecutiveBarrierKernel<<<1, kGinKernelThreads, 0, stream>>>( + rank, kDelayCycles, dObservation, devComm); + ASSERT_MPI_EQ(hipSuccess, hipStreamSynchronize(stream)); + + AllContextsBarrierObservation observation{}; + ASSERT_MPI_EQ(hipSuccess, + hipMemcpy(&observation, dObservation, sizeof(observation), hipMemcpyDeviceToHost)); + if (rank == 0) { + EXPECT_GE(observation.secondBarrierCycles, kDelayCycles / 4) + << "second AllContexts barrier returned before delayed peer arrival"; + EXPECT_EQ(observation.ctxPeerSignals[0], 2u); + EXPECT_EQ(observation.ctxPeerSignals[1], 2u); +#if NCCL_GIN_ANVIL_SDMA_ENABLE + EXPECT_NE(observation.ctxSignalPtrs[0], observation.ctxSignalPtrs[1]) + << "logical contexts must not alias the same signal cell"; +#endif + } + MPI_Barrier(MPI_COMM_WORLD); +} + // Same barrier as barrier2RanksKernel but over the world team, so every rank // has 3 peers instead of 1. __global__ void barrier4RanksKernel(int iters, struct ncclDevComm devComm) { 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 4ab8c21ee1f3..66031c9c6910 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 @@ -127,6 +127,7 @@ { "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": "Barrier_AllContextsConsecutiveSignalsDoNotAlias_SingleNode", "description": "AllContexts barriers keep per-context Anvil signal stripes", "test_filter": "GinMPIDeviceTests.Barrier_AllContextsConsecutiveSignalsDoNotAlias_SingleNode" }, { "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