Reduce RCCL kernel count as we don't pass first coll in argument

src/collectives/device/common.h

@@ -294,7 +294,6 @@ __global__ void NCCL_KERN_NAME(coll, op, dtype)(struct ncclDevComm* comm) { \
   IMPL_COLL_FUNC(coll##LL, op, ncclFunc, dtype, ctype) \
   IMPL_COLL_FUNC(coll##LL128, op, ncclFunc, dtype, ctype) \
   IMPL_COLL_FUNC(coll, op, ncclFunc, dtype, ctype) \
-  IMPL_COLL_KERN_##op(coll##LL, op, ncclFunc, dtype, ctype, FUNC_INDEX(ncclColl, ncclOp, ncclType, al, NCCL_PROTO_LL)) \
 
 #define IMPL_COLL3(coll, op, ncclFunc, dtype, ctype, ncclColl, ncclOp, ncclType) \
   IMPL_COLL4(coll##Tree, op, ncclFunc, dtype, ctype, ncclColl, ncclOp, ncclType, NCCL_ALGO_TREE) \

src/enqueue.cc

@@ -58,15 +58,7 @@
 
 typedef void(*ncclKern_t)(struct ncclDevComm*);
 // Must be consistent with the ncclFuncSet enum
-static ncclKern_t const ncclKerns[4+NCCL_NUM_FUNCTIONS*ncclNumOps*ncclNumTypes*NCCL_NUM_ALGORITHMS*NCCL_NUM_PROTOCOLS] = {
-  NCCL_FUNCS2B(ncclBroadcast),
-  NCCL_FUNCS2A(ncclReduce),
-  NCCL_FUNCS2B(ncclAllGather),
-  NCCL_FUNCS2A(ncclReduceScatter),
-  NCCL_FUNCS2A(ncclAllReduce),
-  NCCL_KERN_NAME(ncclGather, copy, i8),
-  NCCL_KERN_NAME(ncclScatter, copy, i8),
-  NCCL_KERN_NAME(ncclAllToAll, copy, i8),
+static ncclKern_t const ncclKerns[1] = {
   NCCL_KERN_NAME(ncclSendRecv, copy, i8)
 };
 
@@ -128,7 +120,7 @@ ncclResult_t setupLaunch(struct ncclComm* comm, hipLaunchParams* params) {
   struct ncclColl* coll = comm->channels[0].collectives+comm->channels[0].collStart;
 
   comm->args = comm->devComm;
-  params->func = (void *)ncclKerns[coll->funcIndex];
+  params->func = (void *)ncclKerns[0];
   return ncclSuccess;
 }