Plugin version v7

removed old v4 version

Author: bureddy

include/net.h

@@ -22,10 +22,8 @@ typedef enum {NCCL_INIT=1, NCCL_COLL=2, NCCL_P2P=4, NCCL_SHM=8, NCCL_NET=16, NCC
 
 typedef void (*ncclDebugLogger_t)(ncclDebugLogLevel level, unsigned long flags, const char *file, int line, const char *fmt, ...);
 
+#include "net_v7.h"
 #include "net_v6.h"
 #include "net_v5.h"
-#include "net_v4.h"
-
-#define NCCL_PLUGIN_SYMBOL ncclNetPlugin_v6
 
 #endif // end include guard

include/net_device.h

@@ -0,0 +1,31 @@
+/*************************************************************************
+ * Copyright (c) 2023-2023, NVIDIA CORPORATION. All rights reserved.
+ *
+ * See LICENSE.txt for license information
+ ************************************************************************/
+
+#ifndef NET_DEVICE_H_
+#define NET_DEVICE_H_
+
+#include "net_device.h"
+
+#define NCCL_NET_DEVICE_INVALID_VERSION      0x0
+#define NCCL_NET_MTU_SIZE                    4096
+
+// Arbitrary version number - A given NCCL build will only be compatible with a single device networking plugin
+// version. NCCL will check the supplied version number from net->getProperties() and compare to its internal version.
+#define NCCL_NET_DEVICE_UNPACK_VERSION 0x7
+
+typedef enum {NCCL_NET_DEVICE_HOST=0, NCCL_NET_DEVICE_UNPACK=1} ncclNetDeviceType;
+
+typedef struct {
+  ncclNetDeviceType netDeviceType; // Network offload type
+  int netDeviceVersion;            // Version number for network offload
+  void* handle;
+  size_t size;
+  int needsProxyProgress;
+} ncclNetDeviceHandle_v7_t;
+
+typedef ncclNetDeviceHandle_v7_t ncclNetDeviceHandle_t;
+
+#endif

include/net_v6.h

@@ -18,8 +18,6 @@ typedef struct {
   int maxRecvs;   // Maximum number of grouped receives.
 }ncclNetProperties_v6_t;
 
-typedef ncclNetProperties_v6_t ncclNetProperties_t;
-
 typedef struct {
   // Name of the network (mainly for logs)
   const char* name;

include/net_v4.h include/net_v7.hinclude/net_v4.h renamed to include/net_v7.h

@@ -1,22 +1,29 @@
 /*
- * Copyright (c) 2017-2022, NVIDIA CORPORATION. All rights reserved.
+ * Copyright (c) 2017-2023, NVIDIA CORPORATION. All rights reserved.
  */
 
-#ifndef NCCL_NET_V4_H_
-#define NCCL_NET_V4_H_
+#ifndef NCCL_NET_V7_H_
+#define NCCL_NET_V7_H_
+
+#include "net_device.h"
 
 typedef struct {
-  char* name;     // Used mostly for logging.
-  char* pciPath;  // Path to the PCI device in /sys.
-  uint64_t guid;  // Unique identifier for the NIC chip. Important for
-                  // cards with multiple PCI functions (Physical or virtual).
-  int ptrSupport; // NCCL_PTR_HOST or NCCL_PTR_HOST|NCCL_PTR_CUDA
-  int speed;      // Port speed in Mbps.
-  int port;       // Port number.
-  int maxComms;   // Maximum number of comms we can create
-} ncclNetProperties_v4_t;
+  char* name;                      // Used mostly for logging.
+  char* pciPath;                   // Path to the PCI device in /sys.
+  uint64_t guid;                   // Unique identifier for the NIC chip. Important for
+                                   // cards with multiple PCI functions (Physical or virtual).
+  int ptrSupport;                  // [NCCL_PTR_HOST|NCCL_PTR_CUDA|NCCL_PTR_DMABUF]
+  int speed;                       // Port speed in Mbps.
+  int port;                        // Port number.
+  float latency;                   // Network latency
+  int maxComms;                    // Maximum number of comms we can create
+  int maxRecvs;                    // Maximum number of grouped receives.
+  ncclNetDeviceType netDeviceType; // Network offload type
+  int netDeviceVersion;            // Version number for network offload
+} ncclNetProperties_v7_t;
+
+typedef ncclNetProperties_v7_t ncclNetProperties_t;
 
-// v4 struct for backwards compatibility
 typedef struct {
   // Name of the network (mainly for logs)
   const char* name;
@@ -25,37 +32,54 @@ typedef struct {
   // Return the number of adapters.
   ncclResult_t (*devices)(int* ndev);
   // Get various device properties.
-  ncclResult_t (*getProperties)(int dev, ncclNetProperties_v4_t* props);
+  ncclResult_t (*getProperties)(int dev, ncclNetProperties_v7_t* props);
   // Create a receiving object and provide a handle to connect to it. The
   // handle can be up to NCCL_NET_HANDLE_MAXSIZE bytes and will be exchanged
   // between ranks to create a connection.
   ncclResult_t (*listen)(int dev, void* handle, void** listenComm);
   // Connect to a handle and return a sending comm object for that peer.
-  ncclResult_t (*connect)(int dev, void* handle, void** sendComm);
-  // Finalize connection establishment after remote peer has called connectHandle
-  ncclResult_t (*accept)(void* listenComm, void** recvComm);
+  // This call must not block for the connection to be established, and instead
+  // should return successfully with sendComm == NULL with the expectation that
+  // it will be called again until sendComm != NULL.
+  // If *sendDevComm points to a valid object, then NCCL is requesting device offload for this connection
+  ncclResult_t (*connect)(int dev, void* handle, void** sendComm, ncclNetDeviceHandle_v7_t** sendDevComm);
+  // Finalize connection establishment after remote peer has called connect.
+  // This call must not block for the connection to be established, and instead
+  // should return successfully with recvComm == NULL with the expectation that
+  // it will be called again until recvComm != NULL.
+  // If *recvDevComm points to a valid object, then NCCL is requesting device offload for this connection
+  ncclResult_t (*accept)(void* listenComm, void** recvComm, ncclNetDeviceHandle_v7_t** recvDevComm);
   // Register/Deregister memory. Comm can be either a sendComm or a recvComm.
   // Type is either NCCL_PTR_HOST or NCCL_PTR_CUDA.
   ncclResult_t (*regMr)(void* comm, void* data, int size, int type, void** mhandle);
+  /* DMA-BUF support */
+  ncclResult_t (*regMrDmaBuf)(void* comm, void* data, size_t size, int type, uint64_t offset, int fd, void** mhandle);
   ncclResult_t (*deregMr)(void* comm, void* mhandle);
   // Asynchronous send to a peer.
   // May return request == NULL if the call cannot be performed (or would block)
-  ncclResult_t (*isend)(void* sendComm, void* data, int size, void* mhandle, void** request);
+  ncclResult_t (*isend)(void* sendComm, void* data, int size, int tag, void* mhandle, void** request);
   // Asynchronous recv from a peer.
   // May return request == NULL if the call cannot be performed (or would block)
-  ncclResult_t (*irecv)(void* recvComm, void* data, int size, void* mhandle, void** request);
+  ncclResult_t (*irecv)(void* recvComm, int n, void** data, int* sizes, int* tags, void** mhandles, void** request);
   // Perform a flush/fence to make sure all data received with NCCL_PTR_CUDA is
   // visible to the GPU
-  ncclResult_t (*iflush)(void* recvComm, void* data, int size, void* mhandle, void** request);
+  ncclResult_t (*iflush)(void* recvComm, int n, void** data, int* sizes, void** mhandles, void** request);
   // Test whether a request is complete. If size is not NULL, it returns the
   // number of bytes sent/received.
-  ncclResult_t (*test)(void* request, int* done, int* size);
+  ncclResult_t (*test)(void* request, int* done, int* sizes);
   // Close and free send/recv comm objects
   ncclResult_t (*closeSend)(void* sendComm);
   ncclResult_t (*closeRecv)(void* recvComm);
   ncclResult_t (*closeListen)(void* listenComm);
-} ncclNet_v4_t;
 
+  // Copy the given mhandle to a dptr in a format usable by this plugin's device code
+  ncclResult_t (*getDeviceMr)(void* comm, void* mhandle, void** dptr_mhandle);
+
+  // Notify the plugin that a recv has completed by the device
+  ncclResult_t (*irecvConsumed)(void* recvComm, int n, void* request);
+} ncclNet_v7_t;
+
+// v7 struct for backwards compatibility
 typedef struct {
   // Name of the collective network (mainly for logs)
   const char* name;
@@ -65,7 +89,7 @@ typedef struct {
   // If ndev returns 0, all other functions might be set to NULL.
   ncclResult_t (*devices)(int* ndev);
   // Get various device properties.
-  ncclResult_t (*getProperties)(int dev, ncclNetProperties_v4_t* props);
+  ncclResult_t (*getProperties)(int dev, ncclNetProperties_v7_t* props);
   // Create a receiving object and provide a handle to connect to it. The
   // handle can be up to NCCL_NET_HANDLE_MAXSIZE bytes and will be exchanged
   // between ranks to create connections.
@@ -78,6 +102,8 @@ typedef struct {
   ncclResult_t (*reduceSupport)(ncclDataType_t dataType, ncclRedOp_t redOp, int* supported);
   // Register/Deregister memory. Type is either NCCL_PTR_HOST or NCCL_PTR_CUDA.
   ncclResult_t (*regMr)(void* collComm, void* data, int size, int type, void** mhandle);
+  /* DMA-BUF support */
+  ncclResult_t (*regMrDmaBuf)(void* collComm, void* data, size_t size, int type, uint64_t offset, int fd, void** mhandle);
   ncclResult_t (*deregMr)(void* collComm, void* mhandle);
   // Performs an asynchronous allreduce operation on the collective group.
   // May return request == NULL if the call cannot be performed (or would block).
@@ -92,6 +118,6 @@ typedef struct {
   // Close and free collective comm objects
   ncclResult_t (*closeColl)(void* collComm);
   ncclResult_t (*closeListen)(void* listenComm);
-} ncclCollNet_v4_t;
+} ncclCollNet_v7_t;
 
 #endif // end include guard

src/ib_plugin.c

@@ -92,6 +92,24 @@ ncclResult_t ncclIbGetProperties(int dev, ncclNetProperties_t* props)
   return nccl_p2p_ib_get_properties(ncclIbDevs, dev, props);
 }
 
+ncclResult_t ncclIbGetProperties_v6(int dev, ncclNetProperties_v6_t* props_v6)
+{
+  ncclNetProperties_t props;
+  ncclResult_t ret = nccl_p2p_ib_get_properties(ncclIbDevs, dev, &props);
+  if (ret != ncclSuccess) return ret;
+  props_v6->name = props.name;
+  props_v6->pciPath = props.pciPath;
+  props_v6->guid = props.guid;
+  props_v6->ptrSupport = props.ptrSupport;
+  props_v6->speed = props.speed;
+  props_v6->latency = props.latency;
+  props_v6->port = props.port;
+  props_v6->maxComms = props.maxComms;
+  props_v6->maxRecvs = props.maxRecvs;
+
+  return ncclSuccess;
+}
+
 static ncclResult_t GetSocketAddr(union ncclSocketAddress* addr) {
   memcpy(addr, &ncclIbIfAddr, sizeof(*addr));
   return ncclSuccess;
@@ -345,7 +363,7 @@ ncclResult_t ncclIbListen(int dev, void* opaqueHandle, void** listenComm) {
   return ncclSuccess;
 }
 
-ncclResult_t ncclIbConnect(int dev, void* opaqueHandle, void** sendComm) {
+ncclResult_t ncclIbConnect(int dev, void* opaqueHandle, void** sendComm, ncclNetDeviceHandle_t** sendDevComm) {
   struct ncclIbHandle* handle = (struct ncclIbHandle*) opaqueHandle;
   enum ncclSocketState conState;
   struct ncclIbCommStage* stage = &handle->stage;
@@ -461,9 +479,14 @@ ncclResult_t ncclIbConnect(int dev, void* opaqueHandle, void** sendComm) {
   return ncclSuccess;
 }
 
+ncclResult_t ncclIbConnect_v6(int dev, void* opaqueHandle, void** sendComm) {
+  ncclNetDeviceHandle_v7_t* handle = NULL;
+  return ncclIbConnect(dev, opaqueHandle, sendComm, &handle);
+}
+
 NCCL_PARAM(IbGdrFlushDisable, "GDR_FLUSH_DISABLE", 0);
 
-ncclResult_t ncclIbAccept(void* listenComm, void** recvComm) {
+ncclResult_t ncclIbAccept(void* listenComm, void** recvComm,  ncclNetDeviceHandle_t** recvDevComm) {
   struct ncclIbListenComm* lComm = (struct ncclIbListenComm*)listenComm;
   struct ncclIbCommStage* stage = &lComm->stage;
   struct ncclIbRecvComm* rComm = (struct ncclIbRecvComm*)stage->comm;
@@ -604,6 +627,10 @@ ncclResult_t ncclIbAccept(void* listenComm, void** recvComm) {
   stage->buffer = NULL;
   return ncclSuccess;
 }
+ncclResult_t ncclIbAccept_v6(void* listenComm, void** recvComm) {
+  ncclNetDeviceHandle_v7_t* handle = NULL;
+  return ncclIbAccept(listenComm, recvComm, &handle);
+}
 
 ncclResult_t ncclIbGetRequest(struct ncclIbVerbs* verbs, struct ncclIbRequest** req) {
   for (int i=0; i<MAX_REQUESTS; i++) {
@@ -1117,8 +1144,8 @@ ncclResult_t ncclIbCloseListen(void* listenComm) {
   return ncclSuccess;
 }
 
-const ncclNet_v6_t ibPlugin_v6 = {
-  .name = "IBext",
+const ncclNet_v7_t ibPlugin_v7 = {
+  .name = "IBext_v7",
   .init = ncclIbInit,
   .devices = ncclIbDevices,
   .getProperties = ncclIbGetProperties,
@@ -1135,17 +1162,20 @@ const ncclNet_v6_t ibPlugin_v6 = {
   .closeSend = ncclIbCloseSend,
   .closeRecv = ncclIbCloseRecv,
   .closeListen = ncclIbCloseListen,
+  NULL /* getDeviceMr */,
+  NULL /* irecvConsumed */
 };
 
-const ncclNet_v5_t ibPlugin_v5 = {
-  .name = "IBext",
+const ncclNet_v6_t ibPlugin_v6 = {
+  .name = "IBext_v6",
   .init = ncclIbInit,
   .devices = ncclIbDevices,
-  .getProperties = ncclIbGetProperties,
+  .getProperties = ncclIbGetProperties_v6,
   .listen = ncclIbListen,
-  .connect = ncclIbConnect,
-  .accept = ncclIbAccept,
+  .connect = ncclIbConnect_v6,
+  .accept = ncclIbAccept_v6,
   .regMr = ncclIbRegMr,
+  .regMrDmaBuf = ncclIbRegMrDmaBuf,
   .deregMr = ncclIbDeregMr,
   .isend = ncclIbIsend,
   .irecv = ncclIbIrecv,
@@ -1156,58 +1186,19 @@ const ncclNet_v5_t ibPlugin_v5 = {
   .closeListen = ncclIbCloseListen,
 };
 
-static ncclResult_t ncclIbGetProperties_v4(int dev, ncclNetProperties_v4_t* props) {
-  ncclNetProperties_v6_t props_v6;
-  ncclResult_t ret = ncclIbGetProperties(dev, &props_v6);
-  if (ret != ncclSuccess) return ret;
-  props->name = props_v6.name;
-  props->pciPath = props_v6.pciPath;
-  props->guid = props_v6.guid;
-  props->ptrSupport = props_v6.ptrSupport;
-  props->speed = props_v6.speed;
-  props->port = props_v6.port;
-  props->maxComms = props_v6.maxComms;
-  return ncclSuccess;
-}
-static ncclResult_t ncclIbIsend_v4(void *sendComm, void* data, int size, void *mhandle, void** request) {
-  return ncclIbIsend(sendComm, data, size, 0, mhandle, request);
-}
-static ncclResult_t ncclIbIrecv_v4(void* recvComm, void* data, int size, void* mhandle, void** request) {
-  int tag = 0;
-  return ncclIbIrecv(recvComm, 1, &data, &size, &tag, &mhandle, request);
-}
-static ncclResult_t ncclIbIflush_v4(void* recvComm, void* data, int size, void* mhandle, void** request) {
-  return ncclIbIflush(recvComm, 1, &data, &size, &mhandle, request);
-}
-static ncclResult_t ncclIbConnect_v4(int dev, void* handle, void** sendComm) {
-  ncclResult_t ret;
-  do {
-    ret = ncclIbConnect(dev, handle, sendComm);
-  } while (ret == ncclSuccess && *sendComm == NULL);
-  return ret;
-}
-static ncclResult_t ncclIbAccept_v4(void* listenComm, void** recvComm) {
-  ncclResult_t ret;
-  do {
-    ret = ncclIbAccept(listenComm, recvComm);
-  } while (ret == ncclSuccess && *recvComm == NULL);
-  return ret;
-}
-
-
-const ncclNet_v4_t ibPlugin_v4 = {
-  .name = "IBext",
+const ncclNet_v5_t ibPlugin_v5 = {
+  .name = "IBext_v5",
   .init = ncclIbInit,
   .devices = ncclIbDevices,
-  .getProperties = ncclIbGetProperties_v4,
+  .getProperties = ncclIbGetProperties_v6,
   .listen = ncclIbListen,
-  .connect = ncclIbConnect_v4,
-  .accept = ncclIbAccept_v4,
+  .connect = ncclIbConnect_v6,
+  .accept = ncclIbAccept_v6,
   .regMr = ncclIbRegMr,
   .deregMr = ncclIbDeregMr,
-  .isend = ncclIbIsend_v4,
-  .irecv = ncclIbIrecv_v4,
-  .iflush = ncclIbIflush_v4,
+  .isend = ncclIbIsend,
+  .irecv = ncclIbIrecv,
+  .iflush = ncclIbIflush,
   .test = ncclIbTest,
   .closeSend = ncclIbCloseSend,
   .closeRecv = ncclIbCloseRecv,