Skip to content
Merged
Show file tree
Hide file tree
Changes from all 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
9 changes: 8 additions & 1 deletion projects/rccl/src/init.cc
Original file line number Diff line number Diff line change
Expand Up @@ -1626,7 +1626,14 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, struct ncclComm* p
}
}

comm->topo->skipPresetTopoMatching = !uniformRanksPerHost(comm, nranks);
{
bool nonUniformRanks = !uniformRanksPerHost(comm, nranks);
bool isGfx1250 = comm->topo->nodes[GPU].count > 0 && IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx1250");
comm->topo->skipPresetTopoMatching = nonUniformRanks || isGfx1250;
if (comm->topo->skipPresetTopoMatching) {
INFO(NCCL_INIT, "Rome model matching disabled %s", isGfx1250 ? "on gfx1250" : "due to non-uniform ranks per host");
}
}

timers[TIMER_INIT_GRAPHS] = clockNano();
// Get rings and trees
Expand Down
41 changes: 41 additions & 0 deletions projects/rccl/test/RcclWrapTests.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -2386,4 +2386,45 @@ TEST(RcclAllReduceDdaDecision, SymEligible_YieldsToSymmetricKernel)
/*symEligible=*/true, /*ceAllReduceAllowed=*/true));
}

// ---------------------------------------------------------------------------
// Tests for skipPresetTopoMatching: gfx1250 skips Rome model matching.
// Runs on real GPU, initializes a communicator, and checks internal state.
// ---------------------------------------------------------------------------

TEST(SkipPresetTopoMatching, Gfx1250_SkipsRomeModelMatching)
{
RUN_ISOLATED_TEST("Gfx1250_SkipsRomeModelMatching", []()
{
int numDevices = 0;
ASSERT_EQ(hipGetDeviceCount(&numDevices), hipSuccess);
if (numDevices < 1) {
GTEST_SKIP() << "No devices available.";
}

// Check if this is gfx1250
hipDeviceProp_t prop{};
ASSERT_EQ(hipGetDeviceProperties(&prop, 0), hipSuccess);
bool isGfx1250 = (strncmp(prop.gcnArchName, "gfx1250", 7) == 0);
if (!isGfx1250) {
GTEST_SKIP() << "Test only applicable on gfx1250 hardware.";
}

ncclComm_t commHandle{};
ncclUniqueId id{};
ASSERT_EQ(ncclGetUniqueId(&id), ncclSuccess);
ASSERT_EQ(hipSetDevice(0), hipSuccess);
ASSERT_EQ(ncclCommInitRank(&commHandle, 1, id, 0), ncclSuccess);

// Cast to internal struct to access topo
ncclComm* comm = commHandle;

// gfx1250 should skip preset topo matching
EXPECT_TRUE(comm->topo->skipPresetTopoMatching);
// Rome model index should be NONE (no preset matched)
EXPECT_EQ(comm->topo->romeTopoModelIdx, RCCL_ROME_TOPO_PRESET_MODEL_IDX_NONE);

ASSERT_EQ(ncclCommDestroy(commHandle), ncclSuccess);
});
}

} // namespace RcclUnitTesting
Loading