Skip to content
Merged
Show file tree
Hide file tree
Changes from 2 commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
22 changes: 19 additions & 3 deletions projects/rccl/src/plugin/gin.cc
Original file line number Diff line number Diff line change
Expand Up @@ -446,11 +446,27 @@ ncclResult_t ncclGinSetDefaultBackend(struct ncclComm* comm, uint64_t globalBitM
if (ginState->finalized) return ncclSuccess;
std::lock_guard<std::mutex> lock(pluginMutex);

// Close ALL non-global backends and compact the array
// Close ALL non-global backends and compact the array.
// Anvil SDMA peers are local GPU device ids (gin_anvil_sdma_create). On more
// than one node the same id is a different GPU, so leave SDMA out before
// connect instead of letting it become backends[0] or fail the whole GIN
// bring-up (AICOMRCCL-3417). Single-node ordering is unchanged. A node is
// SDMA's reach: one host on MI300X, one MNNVL (UALoE) pod on gfx1250.
for (int i = 0; i < initialNumActiveBackends; i++) {
struct ncclGinBackendState* backend = &ginState->backends[i];
if (!(BIT(backend->ginType) & globalBitMask)) {
WARN("GIN backend type %d not supported on all ranks, removing from active backends", backend->ginType);
#ifdef ENABLE_ROCSHMEM_GIN
const bool anvilMultiNode = comm->nNodes > 1 && backend->ginType == NCCL_GIN_TYPE_ANVIL_SDMA;
Comment thread
dlamd1dai marked this conversation as resolved.
Resolved
Comment thread
dlamd1dai marked this conversation as resolved.
Resolved
#else
const bool anvilMultiNode = false;
#endif
if (anvilMultiNode || !(BIT(backend->ginType) & globalBitMask)) {
Comment thread
dlamd1dai marked this conversation as resolved.
Resolved
if (anvilMultiNode) {
INFO(NCCL_INIT | NCCL_NET,
"GIN backend ANVIL_SDMA is single-node only; removing it from active backends (nNodes=%d)",
comm->nNodes);
} else {
WARN("GIN backend type %d not supported on all ranks, removing from active backends", backend->ginType);
}
NCCLCHECK(ncclGinPluginFinalize(comm, backend->pluginIndex, backend->ginInstance));
} else {
ginState->backends[currentNumActiveBackends++] = *backend;
Expand Down
64 changes: 63 additions & 1 deletion projects/rccl/test/transport/GinDeviceMPITests.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -135,6 +135,23 @@ std::string anvilSingleNodeReason() {
return "";
}

// AICOMRCCL-3417. The rest of this file requires NCCL_GIN_TYPE. This case is
// the opposite: unset, so Anvil SDMA and PROXY are both eligible and SDMA
// would otherwise sort ahead of PROXY on a multi-node communicator.
std::string mixedBackendTwoNodeSkipReason() {
std::string local = ginEnvDisabledReason();
if (local.empty()) local = cuMemReason();
if (local.empty()) {
const char* ginType = std::getenv("NCCL_GIN_TYPE");
if (ginType && ginType[0])
local = "AICOMRCCL-3417 requires NCCL_GIN_TYPE to be unset";
Comment thread
dlamd1dai marked this conversation as resolved.
Resolved
}
// crossNodeReason splits MPI_COMM_TYPE_SHARED. Every rank must call it.
std::string cross = crossNodeReason();
if (local.empty()) local = std::move(cross);
return mpiCoordinatedSkipReason(!local.empty(), local.empty() ? nullptr : local.c_str());
}

// Unlike anvilSingleNodeReason() above, this applies on every backend: the
// fence-visibility tests pair each rank with one peer and read back the peer's
// write, so they are only meaningful with all ranks on one node.
Expand Down Expand Up @@ -436,7 +453,8 @@ class GinMPIDeviceTests : public MPITestBase {
// Minimal 64-byte put + waitSignal round-trip from rank 0 to rank 1 on the
// caller's comm and stream. Contains ASSERT_MPI_*, so it must be the
// caller's last statement.
void runPutWaitRoundTrip(ncclComm_t comm, hipStream_t stream) {
void runPutWaitRoundTrip(ncclComm_t comm, hipStream_t stream,
ncclGinType_t ginType = NCCL_GIN_TYPE_NONE) {
int rank = -1, nRanks = -1;
ncclCommUserRank(comm, &rank);
ncclCommCount(comm, &nRanks);
Expand Down Expand Up @@ -474,6 +492,7 @@ class GinMPIDeviceTests : public MPITestBase {
ncclDevCommRequirements reqs = defaultGinReqs();
reqs.railGinBarrierCount = 1;
reqs.ginSignalCount = 1;
reqs.ginType = ginType;
ncclDevComm devComm{};
ASSERT_MPI_EQ(ncclSuccess, ncclDevCommCreate(comm, &reqs, &devComm));
auto devCommCleanup = makeScopeGuard([&]() {
Expand Down Expand Up @@ -7867,4 +7886,47 @@ TEST_F(GinProxyNthreadsSplitTests, SplitShare_SecondDevCommPut) {
runPutWaitRoundTrip(child, getActiveStream());
}

// Two ranks, two nodes, NCCL_GIN_TYPE unset (AICOMRCCL-3417). Anvil SDMA must
// not stay active: its factory addresses peers by local device id. PROXY is
// the default, a symmetric window registers, a default DevComm and an explicit
// PROXY DevComm each complete a cross-node put/signal, and an explicit Anvil
// DevComm is rejected. Two gfx1250 hosts in one MNNVL (UALoE) pod are one
// node (comm->nNodes == 1); SDMA stays valid there, so the case skips.
TEST_F(GinMPIDeviceTests, MixedBackends_TwoNodeUnsetType) {
if (auto reason = mixedBackendTwoNodeSkipReason(); !reason.empty()) GTEST_SKIP() << reason;
Comment thread
dlamd1dai marked this conversation as resolved.
Resolved
if (auto reason = mpiCoordinatedSkipReason(
!validateTestPrerequisites(/*min_processes=*/2, /*max_processes=*/2,
/*require_power_of_two=*/false, /*min_nodes=*/2, /*max_nodes=*/2),
"Requires exactly 2 ranks on 2 nodes");
!reason.empty())
GTEST_SKIP() << reason;

ASSERT_EQ(ncclSuccess, createTestCommunicator());
SKIP_IF_GIN_UNSUPPORTED();
ncclComm_t comm = getActiveCommunicator();
if (auto reason = mpiCoordinatedSkipReason(comm->nNodes < 2, "Both hosts are in one MNNVL (UALoE) pod");
!reason.empty())
GTEST_SKIP() << reason;
ASSERT_MPI_EQ(2, comm->nNodes);

ncclCommProperties_t props = NCCL_COMM_PROPERTIES_INITIALIZER;
ASSERT_MPI_EQ(ncclSuccess, ncclCommQueryProperties(comm, &props));
ASSERT_MPI_EQ(NCCL_GIN_TYPE_PROXY, props.ginType);
ASSERT_MPI_EQ(true, props.ginSupport[NCCL_GIN_TYPE_PROXY]);
ASSERT_MPI_EQ(false, props.ginSupport[NCCL_GIN_TYPE_ANVIL_SDMA]);

ncclDevCommRequirements anvilReqs = defaultGinReqs();
anvilReqs.ginType = NCCL_GIN_TYPE_ANVIL_SDMA;
anvilReqs.railGinBarrierCount = 1;
anvilReqs.ginSignalCount = 1;
ncclDevComm anvil{};
ncclResult_t anvilSt = ncclDevCommCreate(comm, &anvilReqs, &anvil);
if (anvilSt == ncclSuccess) (void)ncclDevCommDestroy(comm, &anvil);
ASSERT_MPI_EQ(ncclInternalError, anvilSt);

// Default backend (ginType unset on the DevComm) and an explicit PROXY one.
runPutWaitRoundTrip(comm, getActiveStream());
runPutWaitRoundTrip(comm, getActiveStream(), NCCL_GIN_TYPE_PROXY);
}

Comment thread
dlamd1dai marked this conversation as resolved.
Resolved
#endif // MPI_TESTS_ENABLED
Loading