Added inline ctors to LL and LL128 Primitives that match the Simple signature used in all_gather, and explicitly forward params to the proper ctor. This removes the warning and any ambiguity around the implicit overloading of p2pWork and stepSize_.

Author: corey-derochie-amd

src/device/all_gather.h

@@ -70,7 +70,7 @@ namespace {
       // FanSymmetric<1>, only the first element is ever accessed, so it's fine.
       // coverity[callee_ptr_arith:FALSE]
       Primitives<T, RedOp, FanSymmetric<1>, 0, Proto, 0, isNetOffload> prims
-        (tid, workNthreads, &ring->prev, &ring->next, inputBuf, outputBuf, work->redOpArg, 0, work->connIndex, work->connIndex, work, NULL, isNetOffload ? NCCL_MAX_NET_SIZE : 0);
+        (tid, workNthreads, &ring->prev, &ring->next, inputBuf, outputBuf, work->redOpArg, 0, work->connIndex, work->connIndex, work, nullptr, isNetOffload ? NCCL_MAX_NET_SIZE : 0);
 
 #if defined(ENABLE_NPKIT)
       if (tid == 0) {

src/device/prims_ll.h

@@ -680,6 +680,14 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL, P2p, isNetOffload, Metadata, Pi
     setDataPtrs(inputBuf, outputBuf, e != nullptr ? e->acc : nullptr);
   }
 
+  __forceinline__ __device__ Primitives(
+      int tid, int nthreads, int const *recvPeers, int const *sendPeers,
+      void const *inputBuf, void *outputBuf, uint64_t redOpArg, uint8_t group,
+      uint8_t connIndexRecv, uint8_t connIndexSend, struct ncclDevWorkColl* collWork,
+      struct ncclDevWorkP2p* p2pWork, int stepSize_ = 0, int mode = primsModeDefault
+    ): Primitives(tid, nthreads, recvPeers, sendPeers, inputBuf, outputBuf, redOpArg, group,
+                  connIndexRecv, connIndexSend, collWork) {}
+
   __device__ ~Primitives() {
     // Save steps for the next operation
     if (tid >= nthreads-WARP_SIZE && wid < fan.nrecv())

src/device/prims_ll128.h

@@ -600,6 +600,14 @@ class Primitives<T, RedOp, Fan, Direct, ProtoLL128, P2p, isNetOffload, Metadata,
     setDataPtrs(inputBuf, outputBuf, e != nullptr ? e->acc : nullptr);
   }
 
+  __forceinline__ __device__ Primitives(
+      int tid, int nthreads, int const *recvPeers, int const *sendPeers,
+      void const *inputBuf, void *outputBuf, uint64_t redOpArg, uint8_t group,
+      uint8_t connIndexRecv, uint8_t connIndexSend, struct ncclDevWorkColl* collWork,
+      struct ncclDevWorkP2p* p2pWork, int stepSize_ = 0, int mode = primsModeDefault
+    ): Primitives(tid, nthreads, recvPeers, sendPeers, inputBuf, outputBuf, redOpArg, group,
+                  connIndexRecv, connIndexSend, collWork) {}
+
   __device__ ~Primitives() {
     // Save steps for the next operation
     if (tid >= nthreads-WARP_SIZE && wid < fan.nrecv())