Skip to content

Commit 4e5e81e

Browse files
bureddybureddy
committed
Reduce-scatter and Allgather APIs
1 parent d387b8a commit 4e5e81e

8 files changed

Lines changed: 358 additions & 76 deletions

File tree

‎include/net.h‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -22,6 +22,7 @@ typedef enum {NCCL_INIT=1, NCCL_COLL=2, NCCL_P2P=4, NCCL_SHM=8, NCCL_NET=16, NCC
2222

2323
typedef void (*ncclDebugLogger_t)(ncclDebugLogLevel level, unsigned long flags, const char *file, int line, const char *fmt, ...);
2424

25+
#include "net_v8.h"
2526
#include "net_v7.h"
2627
#include "net_v6.h"
2728
#include "net_v5.h"

‎include/net_v7.h‎

Lines changed: 2 additions & 73 deletions
Original file line numberDiff line numberDiff line change
@@ -5,79 +5,8 @@
55
#ifndef NCCL_NET_V7_H_
66
#define NCCL_NET_V7_H_
77

8-
#include "net_device.h"
9-
10-
typedef struct {
11-
char* name; // Used mostly for logging.
12-
char* pciPath; // Path to the PCI device in /sys.
13-
uint64_t guid; // Unique identifier for the NIC chip. Important for
14-
// cards with multiple PCI functions (Physical or virtual).
15-
int ptrSupport; // [NCCL_PTR_HOST|NCCL_PTR_CUDA|NCCL_PTR_DMABUF]
16-
int speed; // Port speed in Mbps.
17-
int port; // Port number.
18-
float latency; // Network latency
19-
int maxComms; // Maximum number of comms we can create
20-
int maxRecvs; // Maximum number of grouped receives.
21-
ncclNetDeviceType netDeviceType; // Network offload type
22-
int netDeviceVersion; // Version number for network offload
23-
} ncclNetProperties_v7_t;
24-
25-
typedef ncclNetProperties_v7_t ncclNetProperties_t;
26-
27-
typedef struct {
28-
// Name of the network (mainly for logs)
29-
const char* name;
30-
// Initialize the network.
31-
ncclResult_t (*init)(ncclDebugLogger_t logFunction);
32-
// Return the number of adapters.
33-
ncclResult_t (*devices)(int* ndev);
34-
// Get various device properties.
35-
ncclResult_t (*getProperties)(int dev, ncclNetProperties_v7_t* props);
36-
// Create a receiving object and provide a handle to connect to it. The
37-
// handle can be up to NCCL_NET_HANDLE_MAXSIZE bytes and will be exchanged
38-
// between ranks to create a connection.
39-
ncclResult_t (*listen)(int dev, void* handle, void** listenComm);
40-
// Connect to a handle and return a sending comm object for that peer.
41-
// This call must not block for the connection to be established, and instead
42-
// should return successfully with sendComm == NULL with the expectation that
43-
// it will be called again until sendComm != NULL.
44-
// If *sendDevComm points to a valid object, then NCCL is requesting device offload for this connection
45-
ncclResult_t (*connect)(int dev, void* handle, void** sendComm, ncclNetDeviceHandle_v7_t** sendDevComm);
46-
// Finalize connection establishment after remote peer has called connect.
47-
// This call must not block for the connection to be established, and instead
48-
// should return successfully with recvComm == NULL with the expectation that
49-
// it will be called again until recvComm != NULL.
50-
// If *recvDevComm points to a valid object, then NCCL is requesting device offload for this connection
51-
ncclResult_t (*accept)(void* listenComm, void** recvComm, ncclNetDeviceHandle_v7_t** recvDevComm);
52-
// Register/Deregister memory. Comm can be either a sendComm or a recvComm.
53-
// Type is either NCCL_PTR_HOST or NCCL_PTR_CUDA.
54-
ncclResult_t (*regMr)(void* comm, void* data, int size, int type, void** mhandle);
55-
/* DMA-BUF support */
56-
ncclResult_t (*regMrDmaBuf)(void* comm, void* data, size_t size, int type, uint64_t offset, int fd, void** mhandle);
57-
ncclResult_t (*deregMr)(void* comm, void* mhandle);
58-
// Asynchronous send to a peer.
59-
// May return request == NULL if the call cannot be performed (or would block)
60-
ncclResult_t (*isend)(void* sendComm, void* data, int size, int tag, void* mhandle, void** request);
61-
// Asynchronous recv from a peer.
62-
// May return request == NULL if the call cannot be performed (or would block)
63-
ncclResult_t (*irecv)(void* recvComm, int n, void** data, int* sizes, int* tags, void** mhandles, void** request);
64-
// Perform a flush/fence to make sure all data received with NCCL_PTR_CUDA is
65-
// visible to the GPU
66-
ncclResult_t (*iflush)(void* recvComm, int n, void** data, int* sizes, void** mhandles, void** request);
67-
// Test whether a request is complete. If size is not NULL, it returns the
68-
// number of bytes sent/received.
69-
ncclResult_t (*test)(void* request, int* done, int* sizes);
70-
// Close and free send/recv comm objects
71-
ncclResult_t (*closeSend)(void* sendComm);
72-
ncclResult_t (*closeRecv)(void* recvComm);
73-
ncclResult_t (*closeListen)(void* listenComm);
74-
75-
// Copy the given mhandle to a dptr in a format usable by this plugin's device code
76-
ncclResult_t (*getDeviceMr)(void* comm, void* mhandle, void** dptr_mhandle);
77-
78-
// Notify the plugin that a recv has completed by the device
79-
ncclResult_t (*irecvConsumed)(void* recvComm, int n, void* request);
80-
} ncclNet_v7_t;
8+
typedef ncclNetProperties_v8_t ncclNetProperties_v7_t;
9+
typedef ncclNet_v8_t ncclNet_v7_t;
8110

8211
// v7 struct for backwards compatibility
8312
typedef struct {

‎include/net_v8.h‎

Lines changed: 134 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,134 @@
1+
/*
2+
* Copyright (c) 2017-2023, NVIDIA CORPORATION. All rights reserved.
3+
*/
4+
5+
#ifndef NCCL_NET_V8_H_
6+
#define NCCL_NET_V8_H_
7+
#include "net_device.h"
8+
9+
typedef struct {
10+
char* name; // Used mostly for logging.
11+
char* pciPath; // Path to the PCI device in /sys.
12+
uint64_t guid; // Unique identifier for the NIC chip. Important for
13+
// cards with multiple PCI functions (Physical or virtual).
14+
int ptrSupport; // [NCCL_PTR_HOST|NCCL_PTR_CUDA|NCCL_PTR_DMABUF]
15+
int speed; // Port speed in Mbps.
16+
int port; // Port number.
17+
float latency; // Network latency
18+
int maxComms; // Maximum number of comms we can create
19+
int maxRecvs; // Maximum number of grouped receives.
20+
ncclNetDeviceType netDeviceType; // Network offload type
21+
int netDeviceVersion; // Version number for network offload
22+
} ncclNetProperties_v8_t;
23+
24+
typedef ncclNetProperties_v8_t ncclNetProperties_t;
25+
26+
typedef struct {
27+
// Name of the network (mainly for logs)
28+
const char* name;
29+
// Initialize the network.
30+
ncclResult_t (*init)(ncclDebugLogger_t logFunction);
31+
// Return the number of adapters.
32+
ncclResult_t (*devices)(int* ndev);
33+
// Get various device properties.
34+
ncclResult_t (*getProperties)(int dev, ncclNetProperties_v8_t* props);
35+
// Create a receiving object and provide a handle to connect to it. The
36+
// handle can be up to NCCL_NET_HANDLE_MAXSIZE bytes and will be exchanged
37+
// between ranks to create a connection.
38+
ncclResult_t (*listen)(int dev, void* handle, void** listenComm);
39+
// Connect to a handle and return a sending comm object for that peer.
40+
// This call must not block for the connection to be established, and instead
41+
// should return successfully with sendComm == NULL with the expectation that
42+
// it will be called again until sendComm != NULL.
43+
// If *sendDevComm points to a valid object, then NCCL is requesting device offload for this connection
44+
ncclResult_t (*connect)(int dev, void* handle, void** sendComm, ncclNetDeviceHandle_v7_t** sendDevComm);
45+
// Finalize connection establishment after remote peer has called connect.
46+
// This call must not block for the connection to be established, and instead
47+
// should return successfully with recvComm == NULL with the expectation that
48+
// it will be called again until recvComm != NULL.
49+
// If *recvDevComm points to a valid object, then NCCL is requesting device offload for this connection
50+
ncclResult_t (*accept)(void* listenComm, void** recvComm, ncclNetDeviceHandle_v7_t** recvDevComm);
51+
// Register/Deregister memory. Comm can be either a sendComm or a recvComm.
52+
// Type is either NCCL_PTR_HOST or NCCL_PTR_CUDA.
53+
ncclResult_t (*regMr)(void* comm, void* data, int size, int type, void** mhandle);
54+
/* DMA-BUF support */
55+
ncclResult_t (*regMrDmaBuf)(void* comm, void* data, size_t size, int type, uint64_t offset, int fd, void** mhandle);
56+
ncclResult_t (*deregMr)(void* comm, void* mhandle);
57+
// Asynchronous send to a peer.
58+
// May return request == NULL if the call cannot be performed (or would block)
59+
ncclResult_t (*isend)(void* sendComm, void* data, int size, int tag, void* mhandle, void** request);
60+
// Asynchronous recv from a peer.
61+
// May return request == NULL if the call cannot be performed (or would block)
62+
ncclResult_t (*irecv)(void* recvComm, int n, void** data, int* sizes, int* tags, void** mhandles, void** request);
63+
// Perform a flush/fence to make sure all data received with NCCL_PTR_CUDA is
64+
// visible to the GPU
65+
ncclResult_t (*iflush)(void* recvComm, int n, void** data, int* sizes, void** mhandles, void** request);
66+
// Test whether a request is complete. If size is not NULL, it returns the
67+
// number of bytes sent/received.
68+
ncclResult_t (*test)(void* request, int* done, int* sizes);
69+
// Close and free send/recv comm objects
70+
ncclResult_t (*closeSend)(void* sendComm);
71+
ncclResult_t (*closeRecv)(void* recvComm);
72+
ncclResult_t (*closeListen)(void* listenComm);
73+
74+
// Copy the given mhandle to a dptr in a format usable by this plugin's device code
75+
ncclResult_t (*getDeviceMr)(void* comm, void* mhandle, void** dptr_mhandle);
76+
77+
// Notify the plugin that a recv has completed by the device
78+
ncclResult_t (*irecvConsumed)(void* recvComm, int n, void* request);
79+
} ncclNet_v8_t;
80+
81+
typedef struct {
82+
void* mhandle;
83+
void* address;
84+
uint32_t size;
85+
} ncclNetSGE_v8_t;
86+
87+
typedef struct {
88+
// Name of the collective network (mainly for logs)
89+
const char* name;
90+
// Initialize the collective network.
91+
ncclResult_t (*init)(ncclDebugLogger_t logFunction);
92+
// Return the number of adapters capable of doing collective operations.
93+
// If ndev returns 0, all other functions might be set to NULL.
94+
ncclResult_t (*devices)(int* ndev);
95+
// Get various device properties.
96+
ncclResult_t (*getProperties)(int dev, ncclNetProperties_v8_t* props);
97+
// Create a receiving object and provide a handle to connect to it. The
98+
// handle can be up to NCCL_NET_HANDLE_MAXSIZE bytes and will be exchanged
99+
// between ranks to create connections.
100+
ncclResult_t (*listen)(int dev, void* handle, void** listenComm);
101+
// Create a group for collective operations. handles have been created
102+
// using listen() above. rank indicates caller's rank in the collective network.
103+
ncclResult_t (*connect)(void* handles[], int nranks, int rank, void* listenComm, void** collComm);
104+
// Returns whether a reduction operation on a data type is supported.
105+
// 1 for supported, 0 otherwise.
106+
ncclResult_t (*reduceSupport)(ncclDataType_t dataType, ncclRedOp_t redOp, int* supported);
107+
// Register/Deregister memory. Type is either NCCL_PTR_HOST or NCCL_PTR_CUDA.
108+
ncclResult_t (*regMr)(void* collComm, void* data, int size, int type, void** mhandle);
109+
/* DMA-BUF support */
110+
ncclResult_t (*regMrDmaBuf)(void* collComm, void* data, size_t size, int type, uint64_t offset, int fd, void** mhandle);
111+
ncclResult_t (*deregMr)(void* collComm, void* mhandle);
112+
// Performs an asynchronous allreduce operation on the collective group.
113+
// May return request == NULL if the call cannot be performed (or would block).
114+
ncclResult_t (*iallreduce)(void* collComm, void* sendData, void* recvData, int count,
115+
ncclDataType_t dataType, ncclRedOp_t redOp, void* sendMhandle, void* recvMhandle, void** request);
116+
ncclResult_t (*iallgather)(void* collComm, void* sendData, int nRecvParts, ncclNetSGE_v8_t* recvParts,
117+
size_t bytesPerRank, size_t windowOffset, size_t windowBytes,
118+
void* sendMhandle, void** request);
119+
ncclResult_t (*ireducescatter)(void* collComm, int nSendParts, ncclNetSGE_v8_t* sendParts, void* recvData,
120+
size_t bytesPerRank, size_t windowOffset, size_t windowBytes,
121+
ncclDataType_t dataType, ncclRedOp_t redOp,
122+
void* recvMhandle, void** request);
123+
// Perform a flush/fence to make sure all data received with NCCL_PTR_CUDA is
124+
// visible to the GPU
125+
ncclResult_t (*iflush)(void* collComm, void* data, int size, void* mhandle, void** request);
126+
// Test whether a request is complete. If size is not NULL, it returns the
127+
// number of bytes sent/received.
128+
ncclResult_t (*test)(void* request, int* done, int* size);
129+
// Close and free collective comm objects
130+
ncclResult_t (*closeColl)(void* collComm);
131+
ncclResult_t (*closeListen)(void* listenComm);
132+
} ncclCollNet_v8_t;
133+
134+
#endif // end include guard

‎src/ib_plugin.c‎

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1144,6 +1144,28 @@ ncclResult_t ncclIbCloseListen(void* listenComm) {
11441144
return ncclSuccess;
11451145
}
11461146

1147+
const ncclNet_v8_t ibPlugin_v8 = {
1148+
.name = "IBext_v8",
1149+
.init = ncclIbInit,
1150+
.devices = ncclIbDevices,
1151+
.getProperties = ncclIbGetProperties,
1152+
.listen = ncclIbListen,
1153+
.connect = ncclIbConnect,
1154+
.accept = ncclIbAccept,
1155+
.regMr = ncclIbRegMr,
1156+
.regMrDmaBuf = ncclIbRegMrDmaBuf,
1157+
.deregMr = ncclIbDeregMr,
1158+
.isend = ncclIbIsend,
1159+
.irecv = ncclIbIrecv,
1160+
.iflush = ncclIbIflush,
1161+
.test = ncclIbTest,
1162+
.closeSend = ncclIbCloseSend,
1163+
.closeRecv = ncclIbCloseRecv,
1164+
.closeListen = ncclIbCloseListen,
1165+
NULL /* getDeviceMr */,
1166+
NULL /* irecvConsumed */
1167+
};
1168+
11471169
const ncclNet_v7_t ibPlugin_v7 = {
11481170
.name = "IBext_v7",
11491171
.init = ncclIbInit,

‎src/p2p_plugin.c‎

Lines changed: 19 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -15,14 +15,17 @@
1515
#include "p2p_plugin.h"
1616

1717
#ifdef HAVE_UCX_PLUGIN
18+
extern ncclNet_v8_t ucxPlugin_v8;
1819
extern ncclNet_v7_t ucxPlugin_v7;
1920
extern ncclNet_v6_t ucxPlugin_v6;
2021
extern ncclNet_v5_t ucxPlugin_v5;
22+
extern ncclNet_v8_t ucxRmaPlugin_v8;
2123
extern ncclNet_v7_t ucxRmaPlugin_v7;
2224
extern ncclNet_v6_t ucxRmaPlugin_v6;
2325
extern ncclNet_v5_t ucxRmaPlugin_v5;
2426
#endif
2527

28+
extern ncclNet_v8_t ibPlugin_v8;
2629
extern ncclNet_v7_t ibPlugin_v7;
2730
extern ncclNet_v6_t ibPlugin_v6;
2831
extern ncclNet_v5_t ibPlugin_v5;
@@ -40,10 +43,16 @@ extern int ncclIbRelaxedOrderingEnabled;
4043
NCCL_PARAM(SharpMaxComms, "SHARP_MAX_COMMS", 1);
4144
NCCL_PARAM(IbAdaptiveRouting, "IB_ADAPTIVE_ROUTING", -2);
4245

46+
ncclResult_t pluginInit_v8(ncclDebugLogger_t logFunction);
4347
ncclResult_t pluginInit_v7(ncclDebugLogger_t logFunction);
4448
ncclResult_t pluginInit_v6(ncclDebugLogger_t logFunction);
4549
ncclResult_t pluginInit_v5(ncclDebugLogger_t logFunction);
4650

51+
ncclNet_v8_t ncclNetPlugin_v8 = {
52+
"NCCL RDMA Plugin v8",
53+
pluginInit_v8,
54+
};
55+
4756
ncclNet_v7_t ncclNetPlugin_v7 = {
4857
"NCCL RDMA Plugin v7",
4958
pluginInit_v7,
@@ -85,17 +94,20 @@ static void pluginSetup()
8594
}
8695
switch (p2p_plugin) {
8796
case NCCL_P2P_IB:
97+
ncclNetPlugin_v8 = ibPlugin_v8;
8898
ncclNetPlugin_v7 = ibPlugin_v7;
8999
ncclNetPlugin_v6 = ibPlugin_v6;
90100
ncclNetPlugin_v5 = ibPlugin_v5;
91101
break;
92102
#ifdef HAVE_UCX_PLUGIN
93103
case NCCL_P2P_UCX:
104+
ncclNetPlugin_v8 = ucxPlugin_v8;
94105
ncclNetPlugin_v7 = ucxPlugin_v7;
95106
ncclNetPlugin_v6 = ucxPlugin_v6;
96107
ncclNetPlugin_v5 = ucxPlugin_v5;
97108
break;
98109
case NCCL_P2P_UCX_RMA:
110+
ncclNetPlugin_v8 = ucxRmaPlugin_v8;
99111
ncclNetPlugin_v7 = ucxRmaPlugin_v7;
100112
ncclNetPlugin_v6 = ucxRmaPlugin_v6;
101113
ncclNetPlugin_v5 = ucxRmaPlugin_v5;
@@ -105,6 +117,13 @@ static void pluginSetup()
105117

106118
}
107119

120+
ncclResult_t pluginInit_v8(ncclDebugLogger_t logFunction) {
121+
pluginLogFunction = logFunction;
122+
pluginSetup();
123+
INFO(NCCL_INIT|NCCL_NET, "P2P plugin %s", ncclNetPlugin_v8.name);
124+
return ncclNetPlugin_v8.init(logFunction);
125+
}
126+
108127
ncclResult_t pluginInit_v7(ncclDebugLogger_t logFunction) {
109128
pluginLogFunction = logFunction;
110129
pluginSetup();

0 commit comments

Comments
 (0)