Removing the experimental clique kernel files (#1610)

Author: gilbertlee-amd

CMakeLists.txt

@@ -424,22 +424,6 @@ set(SRC_FILES
   src/proxy.cc
   src/register.cc
   src/transport.cc
-# src/clique/AllReduceCliqueKernel.h
-# src/clique/CliqueCommon.h
-# src/clique/CliqueManager.cc
-# src/clique/CliqueManager.h
-# src/clique/CliqueShmNames.h
-# src/clique/HandleCache.cc
-# src/clique/HandleCache.h
-# src/clique/HandleShm.cc
-# src/clique/HandleShm.h
-# src/clique/Hash.cc
-# src/clique/Hash.h
-# src/clique/MsgQueue.cc
-# src/clique/MsgQueue.h
-# src/clique/SharedMemHelper.h
-# src/clique/ShmObject.cc
-# src/clique/ShmObject.h
   src/device/all_gather.h
   src/device/all_reduce.h
   src/device/alltoall_pivot.h
@@ -888,7 +872,7 @@ else()
   execute_process(
     COMMAND bash "-c" "free | grep -o '[[:digit:]]*' | head -1"
     OUTPUT_VARIABLE memory_max_string)
-  ## memory_max_string holds the free memory in KB  
+  ## memory_max_string holds the free memory in KB
   if (${memory_max_string} MATCHES "^[0-9]+")
     math(EXPR memory_in_gb "${memory_max_string} / (1024 * 1024)") ## KB to GB conversion
   else()

src/clique/AllReduceCliqueKernel.h

@@ -1059,14 +1059,12 @@ template<typename T, typename RedOp>
 struct RunWorkColl<ncclFuncAllReduce, T, RedOp, NCCL_ALGO_RING, NCCL_PROTO_LL128> {
   __device__ __forceinline__ void run(int tid, int nthreads, struct ncclDevWorkColl* work) {
     runRing<T, RedOp, ProtoLL128>(tid, nthreads, work);
-    //LAUNCH_CLIQUE_KERNEL(AllReduceCliqueSplitKernel, RedOp, T, work);
   }
 };
 
 template<typename T, typename RedOp>
 struct RunWorkColl<ncclFuncAllReduce, T, RedOp, NCCL_ALGO_TREE, NCCL_PROTO_LL128> {
   __device__ __forceinline__ void run(int tid, int nthreads, struct ncclDevWorkColl* work) {
     runTreeSplit<T, RedOp, ProtoLL128>(tid, nthreads, work);
-    //LAUNCH_CLIQUE_KERNEL(AllReduceCliqueSplitKernel, RedOp, T, worrk);
   }
 };

src/clique/CliqueCommon.h

@@ -43,7 +43,6 @@
 #include "git_version.h"
 #include "rccl_vars.h"
 #include "hip_rocm_version_info.h"
-//#include "clique/CliqueManager.h"
 //#include <hsa/hsa_ext_amd.h>
 #ifdef ENABLE_MSCCLPP
 #include "mscclpp/mscclpp_nccl.h"
@@ -485,7 +484,6 @@ static ncclResult_t commFree(ncclComm_t comm) {
   return ncclSuccess;
 }
 
-RCCL_PARAM(CliqueIgnoreTopo, "CLIQUE_IGNORE_TOPO", 0);
 RCCL_PARAM(P2pNetDisable, "P2P_NET_DISABLE", 0);
 RCCL_PARAM(PivotAlltoallEnable, "PIVOT_ALLTOALL_ENABLE", 1);
 RCCL_PARAM(LL128ForceEnable, "LL128_FORCE_ENABLE", 0);
@@ -1578,7 +1576,7 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, struct ncclComm* p
   } else {
     NCCLCHECKGOTO(ncclProxyCreate(comm), ret, fail);
   }
-  
+
   timers[TIMER_INIT_CONNECT] = clockNano();
   do { // Build p2p schedule
     int node = comm->node;
@@ -1979,7 +1977,7 @@ static ncclResult_t ncclCommInitRankFunc(struct ncclAsyncJob* job_) {
 	if (rcclParamMscclppForceEnabled()) {
 		comm->mscclppForceEnable = true;
 	} else {
-		comm->mscclppForceEnable = false;	
+		comm->mscclppForceEnable = false;
 	}
       } else {
         WARN("MSCCL++: Cannot enable MSCCL++ on %s architecture", devProp.gcnArchName);
@@ -2410,7 +2408,7 @@ static ncclResult_t commDestroySync(struct ncclAsyncJob* job_) {
     // And keep polling until all graphs referencing us die.
     while (comm->persistentRefs != 0) {
       NCCLCHECKGOTO(ncclCommPollCallbacks(comm, /*waitSome=*/true), ret, fail);
-    }  
+    }
   }
 
   if ((ret = ncclProxyStop(comm)) != ncclSuccess) {

src/clique/CliqueManager.cc

@@ -118,7 +118,7 @@ namespace RcclUnitTesting
         std::vector<bool>           const managedMemList  = {false};
         std::vector<bool>           const useHipGraphList = {false, true};
         std::vector<const char *>   const channelList     = {"84", "112"};
-        bool                        const enableSweep     = false; 
+        bool                        const enableSweep     = false;
         for (auto channel : channelList) {
           setenv("NCCL_MIN_NCHANNELS", channel, 1);
           testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
@@ -149,30 +149,6 @@ namespace RcclUnitTesting
     testBed.Finalize();
   }
 
-  TEST(AllReduce, DISABLED_Clique)
-  {
-    // Set clique env var prior to TestBed
-    setenv("RCCL_ENABLE_CLIQUE", "1", 1);
-
-    TestBed testBed;
-
-    // Configuration
-    std::vector<ncclFunc_t>     const funcTypes       = {ncclCollAllReduce};
-    std::vector<ncclDataType_t> const dataTypes       = testBed.GetAllSupportedDataTypes();
-    std::vector<ncclRedOp_t>    const redOps          = testBed.GetAllSupportedRedOps();
-    std::vector<int>            const roots           = {0};
-    std::vector<int>            const numElements     = {1048576, 1024};
-    std::vector<bool>           const inPlaceList     = {false, true};
-    std::vector<bool>           const managedMemList  = {false};
-    std::vector<bool>           const useHipGraphList = {false, true};
-
-    testBed.RunSimpleSweep(funcTypes, dataTypes, redOps, roots, numElements,
-                           inPlaceList, managedMemList, useHipGraphList);
-    testBed.Finalize();
-
-    unsetenv("RCCL_ENABLE_CLIQUE");
-  }
-
   // This tests using custom pre-mult scalars reductions
   TEST(AllReduce, PreMultScalar)
   {
@@ -245,7 +221,7 @@ namespace RcclUnitTesting
   }
 
   TEST(AllReduce, UserBufferRegistration)
-  {          
+  {
     const int nranks = 8;
     size_t count = 2048;
     std::vector<int> sendBuff(count, 0);
@@ -260,7 +236,7 @@ namespace RcclUnitTesting
   }
 
   TEST(AllReduce, ManagedMemUserBufferRegistration)
-  {          
+  {
     const int nranks = 8;
     size_t count = 2048;
     std::vector<int> sendBuff(count, 0);

src/clique/CliqueManager.h

@@ -3,19 +3,10 @@ RCCL_INSTALL=../../build/release
 EXE=$PWD/HelloRccl
 LDPATH=$LD_LIBRARY_PATH:$RCCL_INSTALL
 
-echo "Single process - With clique-based kernels:"
-RCCL_CLIQUE_ALLREDUCE_BYTE_LIMIT=1073741824 RCCL_FORCE_ENABLE_CLIQUE=1 NCCL_DEBUG=INFO RCCL_ENABLE_CLIQUE=1 LD_LIBRARY_PATH=$LDPATH $EXE 4
-
-echo "Single process - Without clique-based kernels:"
+echo "Single process:"
 NCCL_DEBUG=INFO LD_LIBRARY_PATH=$LDPATH $EXE 4
 
-echo "With clique-based kernels:"
-RCCL_CLIQUE_ALLREDUCE_BYTE_LIMIT=1073741824 RCCL_FORCE_ENABLE_CLIQUE=1 NCCL_DEBUG=INFO RCCL_ENABLE_CLIQUE=1 NCCL_COMM_ID=$HOSTNAME:12345 LD_LIBRARY_PATH=$LDPATH $EXE 4 0 &
-RCCL_CLIQUE_ALLREDUCE_BYTE_LIMIT=1073741824 RCCL_FORCE_ENABLE_CLIQUE=1 NCCL_DEBUG=INFO RCCL_ENABLE_CLIQUE=1 NCCL_COMM_ID=$HOSTNAME:12345 LD_LIBRARY_PATH=$LDPATH $EXE 4 1 &
-RCCL_CLIQUE_ALLREDUCE_BYTE_LIMIT=1073741824 RCCL_FORCE_ENABLE_CLIQUE=1 NCCL_DEBUG=INFO RCCL_ENABLE_CLIQUE=1 NCCL_COMM_ID=$HOSTNAME:12345 LD_LIBRARY_PATH=$LDPATH $EXE 4 2 &
-RCCL_CLIQUE_ALLREDUCE_BYTE_LIMIT=1073741824 RCCL_FORCE_ENABLE_CLIQUE=1 NCCL_DEBUG=INFO RCCL_ENABLE_CLIQUE=1 NCCL_COMM_ID=$HOSTNAME:12345 LD_LIBRARY_PATH=$LDPATH $EXE 4 3
-
-echo "Without clique-based kernels:"
+echo "Multi-process:"
 NCCL_COMM_ID=$HOSTNAME:12345 LD_LIBRARY_PATH=$LDPATH $EXE 4 0 &
 NCCL_COMM_ID=$HOSTNAME:12345 LD_LIBRARY_PATH=$LDPATH $EXE 4 1 &
 NCCL_COMM_ID=$HOSTNAME:12345 LD_LIBRARY_PATH=$LDPATH $EXE 4 2 &

src/clique/CliqueShmNames.h

@@ -38,7 +38,6 @@ const char* ncclProtoStr[NCCL_NUM_PROTOCOLS] = { "LL", "LL128", "Simple" };
 
 extern NodeModel *node_model;
 
-RCCL_PARAM(CliqueIgnoreTopo, "CLIQUE_IGNORE_TOPO", 0);
 RCCL_PARAM(P2pNetDisable, "P2P_NET_DISABLE", 0);
 RCCL_PARAM(PivotAlltoallEnable, "PIVOT_ALLTOALL_ENABLE", 1);
 RCCL_PARAM(LL128ForceEnable, "LL128_FORCE_ENABLE", 0);