diff --git a/src/collectives/device/common.h b/src/collectives/device/common.h index 7f834ddc4e..da67970e45 100644 --- a/src/collectives/device/common.h +++ b/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) \ diff --git a/src/enqueue.cc b/src/enqueue.cc index 97ccb5b6e3..a5438b5b27 100644 --- a/src/enqueue.cc +++ b/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; }