Revert "Enable Ll128 on gfx90a (#823)" (#829)

This reverts commit 420f8af6a0.

Also increase number of parallel jobs for linking

CMakeLists.txt

@@ -593,7 +593,7 @@ endif()
 
 ## Set RCCL link options
 target_link_options(rccl PRIVATE -fgpu-rdc)               # Required when linking relocatable device code
-target_link_options(rccl PRIVATE -parallel-jobs=12)       # Use multiple threads to link
+target_link_options(rccl PRIVATE -parallel-jobs=16)       # Use multiple threads to link
 if(BUILD_ADDRESS_SANITIZER)
   target_link_options(rccl PRIVATE -fuse-ld=lld)
 endif()

src/collectives/device/common.h

@@ -28,7 +28,7 @@
   { __atomic_store_n((DST), (SRC), __ATOMIC_SEQ_CST); }
 #endif
 
-#ifdef __gfx90a__
+#ifdef ENABLE_LL128
 #define NCCL_FUNC5(func, algo, devredop, type, nullify) \
   MACRO_IF(nullify, nullptr, NCCL_FUNC_NAME(func, algo, LL,     devredop, type)), \
   MACRO_IF(nullify, nullptr, NCCL_FUNC_NAME(func, algo, LL128,  devredop, type)), \
@@ -549,7 +549,7 @@ __forceinline__ __device__ void ncclKernel(
 #ifdef USE_INDIRECT_FUNCTION_CALL
       ncclFuncs[ncclShmem.work.header.funcIndex]();
 #else
-#ifdef __gfx90a__
+#ifdef ENABLE_LL128
       NCCL_CALL_FUNCTIONS<1>(ncclShmem.work.header.funcIndex);
 #else
       NCCL_CALL_FUNCTIONS<0>(ncclShmem.work.header.funcIndex);
@@ -618,7 +618,7 @@ __device__  __attribute__((noinline)) void NCCL_FUNC_NAME(func, algo, proto, dev
 #endif
 
 // Only generate inline kernels for LL
-#ifdef __gfx90a__
+#ifdef ENABLE_LL128
 #define IMPL_COLL4(func, algo, devredop, type) \
   IMPL_COLL_FUNC(func, algo, LL,     devredop, type) \
   IMPL_COLL_FUNC(func, algo, LL128,  devredop, type) \

src/graph/tuning.cc

@@ -451,9 +451,13 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom
     int pEnable = protoEnable[p];
     if (pEnable == 2 && p == NCCL_PROTO_LL128) {
 #if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__)
+#if defined(ENABLE_LL128)
       // Enable LL128 by default only on gfx90a with available tuning table
       pEnable = (graphs[a]->typeInter <= PATH_PXB) && graphs[a]->typeIntra <= PATH_NVL &&
         (comm->topo->nodes[GPU].nodes[0].gpu.gcn == 910 && comm->topo->ll128Enabled) ? 1 : 0;
+#else
+      pEnable = 0;
+#endif
 #else
       // Enable LL128 by default only on Volta/Ampere/Hopper+NVLink. Other cases are not tested and may cause silent data corruption.
       pEnable = 1;