gfx12 Disable ll protocol (#1268)

2024-07-26 08:59:55 -06:00
@@ -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;
@@ -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; c<NCCL_NUM_FUNCTIONS; c++) for (int a=0; a<NCCL_NUM_ALGORITHMS; a++) for (int p=0; p<NCCL_NUM_PROTOCOLS; p++) {
-    // Disable LL protocol on gfx11xx
-    int pEnable = protoEnable[p];
+    // Disable LL protocol on gfx12xx
+    int pEnable = (p == NCCL_PROTO_LL && IsArchMatch(comm->topo->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)
@@ -7,9 +7,10 @@
 #include <gtest/gtest.h>
 #include <rccl/rccl.h>

+#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<ncclComm_t> subComms(numDevices);  
+    std::vector<ncclComm_t> 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<ncclComm_t> 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
@@ -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);
@@ -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();