diff --git a/src/enqueue.cc b/src/enqueue.cc index ef548df111..96db86bd03 100644 --- a/src/enqueue.cc +++ b/src/enqueue.cc @@ -271,7 +271,7 @@ static ncclResult_t computeCollAlignCount(struct ncclInfo* collInfo, size_t* ali if (collInfo->protocol == NCCL_PROTO_SIMPLE) { *alignCount = NCCL_SIMPLE_ALIGNMENT / ncclTypeSize(collInfo->datatype); } else if (collInfo->protocol == NCCL_PROTO_LL128) { - // LL128 alignCount should be same as LL for now. NCCL_LL128_ALIGNMENT_PER_WARP needs review + // LL128 alignCount should be same as LL for now. NCCL_LL128_ALIGNMENT_PER_WARP needs review *alignCount = NCCL_LL_ALIGNMENT_PER_THREAD / ncclTypeSize(collInfo->datatype) * collInfo->nThreads; } else { *alignCount = NCCL_LL_ALIGNMENT_PER_THREAD / ncclTypeSize(collInfo->datatype) * collInfo->nThreads; @@ -591,8 +591,8 @@ static ncclResult_t addP2pToPlan( // 1 is connIndex struct ncclConnInfo* conn = isSendNotRecv ? &comm->channels[channelId].peers[peer]->send[1].conn : &comm->channels[channelId].peers[peer]->recv[1].conn; - // do not use LL on gfx11 - info.protocol = ((conn->buffs[NCCL_PROTO_LL] != nullptr) && bytes <= ncclParamP2pLLThreshold()) ? NCCL_PROTO_LL : NCCL_PROTO_SIMPLE; + // do not use LL on gfx12 + info.protocol = ((conn->buffs[NCCL_PROTO_LL] != nullptr) && bytes <= ncclParamP2pLLThreshold() && !IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx12")) ? NCCL_PROTO_LL : NCCL_PROTO_SIMPLE; int reg = 0; if (info.protocol == NCCL_PROTO_SIMPLE) { @@ -1923,7 +1923,7 @@ static ncclResult_t hostToDevRedOp( int64_t i64; uint64_t u64; half f16; - float f32; + float f32; double f64; #if defined(RCCL_BFLOAT16) hip_bfloat16 bf16; diff --git a/src/graph/tuning.cc b/src/graph/tuning.cc index a67cb17e11..57dd23d272 100644 --- a/src/graph/tuning.cc +++ b/src/graph/tuning.cc @@ -54,7 +54,7 @@ ncclResult_t parseList(const char* str, const char* elems[], int nelems, int* li // Latencies in us, Bandwidths in GB/s // Tree { LL, LL128, Simple } , Ring { LL, LL128, Simple } -static const float baseLat [NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = { +static const float baseLat [NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = { { 12.0, 12.0, 17.0 }, { 12.0, 12.0, 17.0 }, // Tree, Ring { 12.0, 12.0, 17.0 }, { 12.0, 12.0, 17.0 }, // Collnet Direct, Chain { 0, 0, 0 }, { 0, 0, 0 }}; // NVLS, NVLS Tree @@ -372,7 +372,7 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom if (a == NCCL_ALGO_NVLS) bw = std::min(graphs[a]->bwIntra, graphs[a]->bwInter); if (a == NCCL_ALGO_NVLS_TREE) bw = std::min(graphs[a]->bwIntra, nNodes <= 2 ? graphs[a]->bwInter : graphs[a]->bwInter/2); - + // Various model refinements #if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) if (nNodes <= 2) @@ -507,8 +507,8 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom } for (int c=0; ctopo->nodes[GPU].nodes[0].gpu.gcn, "gfx12")) ? 0 : protoEnable[p]; if (pEnable == 2 && p == NCCL_PROTO_LL128) { #if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) #if defined(ENABLE_LL128) diff --git a/test/StandaloneTests.cpp b/test/StandaloneTests.cpp index 887c15204d..cbb73efde5 100644 --- a/test/StandaloneTests.cpp +++ b/test/StandaloneTests.cpp @@ -7,9 +7,10 @@ #include #include +#include "TestBed.hpp" #include "StandaloneUtils.hpp" -namespace RcclUnitTesting +namespace RcclUnitTesting { /** * \brief Verify that each device is assigned to the right rank using ncclCommSplit API. @@ -73,7 +74,7 @@ namespace RcclUnitTesting NCCLCHECK(ncclCommInitAll(comms.data(), numDevices, nullptr)); // Split into new comms (all of the same color) - std::vector subComms(numDevices); + std::vector subComms(numDevices); NCCLCHECK(ncclGroupStart()); for (int localRank = 0; localRank < numDevices; localRank++) NCCLCHECK(ncclCommSplit(comms[localRank], 0, localRank, &subComms[localRank], NULL)); @@ -88,7 +89,7 @@ namespace RcclUnitTesting int subCommRank, subCommNRank; NCCLCHECK(ncclCommUserRank(subComms[i], &subCommRank)); NCCLCHECK(ncclCommCount(subComms[i], &subCommNRank)); - + ASSERT_EQ(originalRank, subCommRank); ASSERT_EQ(originalNRank, subCommNRank); } @@ -117,7 +118,7 @@ namespace RcclUnitTesting NCCLCHECK(ncclCommInitAll(comms.data(), numDevices, nullptr)); // Split into new comms - int numReducedRanks = numDevices / 2; + int numReducedRanks = numDevices / 2; std::vector subComms(numDevices); NCCLCHECK(ncclGroupStart()); for (int localRank = 0; localRank < numDevices; localRank++) @@ -131,12 +132,12 @@ namespace RcclUnitTesting int originalRank, originalNRank; NCCLCHECK(ncclCommUserRank(comms[i], &originalRank)); NCCLCHECK(ncclCommCount(comms[i], &originalNRank)); - + if (i < numReducedRanks) { int subCommRank, subCommNRank; NCCLCHECK(ncclCommUserRank(subComms[i], &subCommRank)); NCCLCHECK(ncclCommCount(subComms[i], &subCommNRank)); - + ASSERT_EQ(originalRank, subCommRank); ASSERT_EQ(subCommNRank, numReducedRanks); } else { @@ -150,12 +151,13 @@ namespace RcclUnitTesting for (auto& comm : comms) NCCLCHECK(ncclCommDestroy(comm)); } - + /** * \brief Verify there is no regression in timing for each protocol [LL, LL128, Simple] * ******************************************************************************************/ TEST(Standalone, RegressionTiming) { + TestBed testBed; // timing using namespace std::chrono; using Clock = std::chrono::high_resolution_clock; @@ -178,7 +180,12 @@ namespace RcclUnitTesting for (auto p : protocolList) { usElapsed = 0; - setenv("NCCL_PROTO", p, 1); + if(testBed.ev.isGfx12) { + setenv("NCCL_PROTO", "Simple", 1); + } else { + setenv("NCCL_PROTO", p, 1); + } + NCCLCHECK(ncclCommInitAll(comms.data(), numRanks, nullptr)); // Prepare CPU data arrays diff --git a/test/common/EnvVars.cpp b/test/common/EnvVars.cpp index c51b7be012..d7fd501891 100644 --- a/test/common/EnvVars.cpp +++ b/test/common/EnvVars.cpp @@ -15,7 +15,7 @@ namespace RcclUnitTesting int const UT_SINGLE_PROCESS = (1<<0); int const UT_MULTI_PROCESS = (1<<1); - int getArchInfo(bool *isRightArch) + int getArchInfo(bool *isRightArch, const char *gfx) { // Prepare parent->child pipe int pipefd[2]; @@ -25,7 +25,7 @@ namespace RcclUnitTesting } pid_t pid = fork(); if (0 == pid) { - bool isGfx94 = false; + bool isGfxTest = false; int dev; hipGetDeviceCount(&dev); for (int deviceId = 0; deviceId < dev; deviceId++) { @@ -34,14 +34,14 @@ namespace RcclUnitTesting hipGetDeviceProperties(&devProp, deviceId); char *gcnArchNameToken = strtok(devProp.gcnArchName, ":"); strcpy(gcn, gcnArchNameToken); - if(std::strncmp("gfx94", gcn, 5) == 0) { - isGfx94 = true; + if(std::strncmp(gfx, gcn, 5) == 0) { + isGfxTest = true; } else { - isGfx94 = false; + isGfxTest = false; break; } } - if (write(pipefd[1], &isGfx94, sizeof(isGfx94)) != sizeof(isGfx94)) return TEST_FAIL; + if (write(pipefd[1], &isGfxTest, sizeof(isGfxTest)) != sizeof(isGfxTest)) return TEST_FAIL; close(pipefd[0]); close(pipefd[1]); exit(EXIT_SUCCESS); @@ -95,7 +95,9 @@ namespace RcclUnitTesting numDetectedGpus = 0; getDeviceCount(&numDetectedGpus); isGfx94 = false; - getArchInfo(&isGfx94); + getArchInfo(&isGfx94, "gfx94"); + isGfx12 = false; + getArchInfo(&isGfx12, "gfx12"); showNames = GetEnvVar("UT_SHOW_NAMES" , 1); minGpus = GetEnvVar("UT_MIN_GPUS" , 2); diff --git a/test/common/EnvVars.hpp b/test/common/EnvVars.hpp index ea122e5971..c20e82cf45 100644 --- a/test/common/EnvVars.hpp +++ b/test/common/EnvVars.hpp @@ -31,6 +31,7 @@ namespace RcclUnitTesting int timeoutUs; // Set timeout for child in microseconds [UT_TIMEOUT_US] bool useMultithreading; // Multi-thread single-process ranks [UT_MULTITHREAD] bool isGfx94; // Detects if architecture is gfx94 + bool isGfx12; // Detects if architecture is gfx12 // Constructor that parses and collects environment variables EnvVars();