Skip to content

Commit 91aa77a

Browse files
MC952-archMC952-arch
authored
Add zero-copy support for P2P operation (#261)
* Add flagcxRegPool for registration management * Implement flagcxProxyMsgRegister/Deregister operations * Add CCLAdaptor memAlloc/Free and commRegister/Deregister APIs * Refactor P2P op execution from op-wise to group-wise * Add flagcxHostSemaphore for P2P op tracking * Extend test/perf with '-R [0/1]' argument for user buffer registration * Add DeviceAdaptor hostGetDevicePointer API
1 parent 2f6f188 commit 91aa77a

48 files changed

Lines changed: 1406 additions & 303 deletions

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

‎flagcx/adaptor/ccl/bootstrap_adaptor.cc‎

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -84,6 +84,26 @@ flagcxResult_t bootstrapAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
8484
return flagcxNotSupported;
8585
}
8686

87+
// TODO: unsupported
88+
flagcxResult_t bootstrapAdaptorMemAlloc(void **ptr, size_t size) {
89+
return flagcxNotSupported;
90+
}
91+
92+
// TODO: unsupported
93+
flagcxResult_t bootstrapAdaptorMemFree(void *ptr) { return flagcxNotSupported; }
94+
95+
// TODO: unsupported
96+
flagcxResult_t bootstrapAdaptorCommRegister(flagcxInnerComm_t comm, void *buff,
97+
size_t size, void **handle) {
98+
return flagcxNotSupported;
99+
}
100+
101+
// TODO: unsupported
102+
flagcxResult_t bootstrapAdaptorCommDeregister(flagcxInnerComm_t comm,
103+
void *handle) {
104+
return flagcxNotSupported;
105+
}
106+
87107
flagcxResult_t bootstrapAdaptorGather(const void *sendbuff, void *recvbuff,
88108
size_t count, flagcxDataType_t datatype,
89109
int root, flagcxInnerComm_t comm,
@@ -212,6 +232,8 @@ struct flagcxCCLAdaptor bootstrapAdaptor = {
212232
bootstrapAdaptorCommResume, bootstrapAdaptorCommSuspend,
213233
bootstrapAdaptorCommCount, bootstrapAdaptorCommCuDevice,
214234
bootstrapAdaptorCommUserRank, bootstrapAdaptorCommGetAsyncError,
235+
bootstrapAdaptorMemAlloc, bootstrapAdaptorMemFree,
236+
bootstrapAdaptorCommRegister, bootstrapAdaptorCommDeregister,
215237
// Communication functions
216238
bootstrapAdaptorReduce, bootstrapAdaptorGather, bootstrapAdaptorScatter,
217239
bootstrapAdaptorBroadcast, bootstrapAdaptorAllReduce,

‎flagcx/adaptor/ccl/cncl_adaptor.cc‎

Lines changed: 21 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -113,6 +113,25 @@ flagcxResult_t cnclAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
113113
return flagcxSuccess;
114114
}
115115

116+
// TODO: unsupported
117+
flagcxResult_t cnclAdaptorMemAlloc(void **ptr, size_t size) {
118+
return flagcxNotSupported;
119+
}
120+
121+
// TODO: unsupported
122+
flagcxResult_t cnclAdaptorMemFree(void *ptr) { return flagcxNotSupported; }
123+
124+
// TODO: unsupported
125+
flagcxResult_t cnclAdaptorCommRegister(flagcxInnerComm_t comm, void *buff,
126+
size_t size, void **handle) {
127+
return flagcxNotSupported;
128+
}
129+
130+
// TODO: unsupported
131+
flagcxResult_t cnclAdaptorCommDeregister(flagcxInnerComm_t comm, void *handle) {
132+
return flagcxNotSupported;
133+
}
134+
116135
flagcxResult_t cnclAdaptorReduce(const void *sendbuff, void *recvbuff,
117136
size_t count, flagcxDataType_t datatype,
118137
flagcxRedOp_t op, int root,
@@ -305,7 +324,8 @@ struct flagcxCCLAdaptor cnclAdaptor = {
305324
cnclAdaptorCommInitRank, cnclAdaptorCommFinalize, cnclAdaptorCommDestroy,
306325
cnclAdaptorCommAbort, cnclAdaptorCommResume, cnclAdaptorCommSuspend,
307326
cnclAdaptorCommCount, cnclAdaptorCommCuDevice, cnclAdaptorCommUserRank,
308-
cnclAdaptorCommGetAsyncError,
327+
cnclAdaptorCommGetAsyncError, cnclAdaptorMemAlloc, cnclAdaptorMemFree,
328+
cnclAdaptorCommRegister, cnclAdaptorCommDeregister,
309329
// Communication functions
310330
cnclAdaptorReduce, cnclAdaptorGather, cnclAdaptorScatter,
311331
cnclAdaptorBroadcast, cnclAdaptorAllReduce, cnclAdaptorReduceScatter,

‎flagcx/adaptor/ccl/dunccl_adaptor.cc‎

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -72,6 +72,26 @@ flagcxResult_t duncclAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
7272
(ncclResult_t *)&asyncError);
7373
}
7474

75+
// TODO: unsupported
76+
flagcxResult_t duncclAdaptorMemAlloc(void **ptr, size_t size) {
77+
return flagcxNotSupported;
78+
}
79+
80+
// TODO: unsupported
81+
flagcxResult_t duncclAdaptorMemFree(void *ptr) { return flagcxNotSupported; }
82+
83+
// TODO: unsupported
84+
flagcxResult_t duncclAdaptorCommRegister(flagcxInnerComm_t comm, void *buff,
85+
size_t size, void **handle) {
86+
return flagcxNotSupported;
87+
}
88+
89+
// TODO: unsupported
90+
flagcxResult_t duncclAdaptorCommDeregister(flagcxInnerComm_t comm,
91+
void *handle) {
92+
return flagcxNotSupported;
93+
}
94+
7595
flagcxResult_t duncclAdaptorReduce(const void *sendbuff, void *recvbuff,
7696
size_t count, flagcxDataType_t datatype,
7797
flagcxRedOp_t op, int root,
@@ -261,6 +281,8 @@ struct flagcxCCLAdaptor duncclAdaptor = {
261281
duncclAdaptorCommDestroy, duncclAdaptorCommAbort, duncclAdaptorCommResume,
262282
duncclAdaptorCommSuspend, duncclAdaptorCommCount, duncclAdaptorCommCuDevice,
263283
duncclAdaptorCommUserRank, duncclAdaptorCommGetAsyncError,
284+
duncclAdaptorMemAlloc, duncclAdaptorMemFree, duncclAdaptorCommRegister,
285+
duncclAdaptorCommDeregister,
264286
// Communication functions
265287
duncclAdaptorReduce, duncclAdaptorGather, duncclAdaptorScatter,
266288
duncclAdaptorBroadcast, duncclAdaptorAllReduce, duncclAdaptorReduceScatter,

‎flagcx/adaptor/ccl/gloo_adaptor.cc‎

Lines changed: 21 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -115,6 +115,25 @@ flagcxResult_t glooAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
115115
return flagcxNotSupported;
116116
}
117117

118+
// TODO: unsupported
119+
flagcxResult_t glooAdaptorMemAlloc(void **ptr, size_t size) {
120+
return flagcxNotSupported;
121+
}
122+
123+
// TODO: unsupported
124+
flagcxResult_t glooAdaptorMemFree(void *ptr) { return flagcxNotSupported; }
125+
126+
// TODO: unsupported
127+
flagcxResult_t glooAdaptorCommRegister(flagcxInnerComm_t comm, void *buff,
128+
size_t size, void **handle) {
129+
return flagcxNotSupported;
130+
}
131+
132+
// TODO: unsupported
133+
flagcxResult_t glooAdaptorCommDeregister(flagcxInnerComm_t comm, void *handle) {
134+
return flagcxNotSupported;
135+
}
136+
118137
flagcxResult_t glooAdaptorReduce(const void *sendbuff, void *recvbuff,
119138
size_t count, flagcxDataType_t datatype,
120139
flagcxRedOp_t op, int root,
@@ -305,7 +324,8 @@ struct flagcxCCLAdaptor glooAdaptor = {
305324
glooAdaptorCommInitRank, glooAdaptorCommFinalize, glooAdaptorCommDestroy,
306325
glooAdaptorCommAbort, glooAdaptorCommResume, glooAdaptorCommSuspend,
307326
glooAdaptorCommCount, glooAdaptorCommCuDevice, glooAdaptorCommUserRank,
308-
glooAdaptorCommGetAsyncError,
327+
glooAdaptorCommGetAsyncError, glooAdaptorMemAlloc, glooAdaptorMemFree,
328+
glooAdaptorCommRegister, glooAdaptorCommDeregister,
309329
// Communication functions
310330
glooAdaptorReduce, glooAdaptorGather, glooAdaptorScatter,
311331
glooAdaptorBroadcast, glooAdaptorAllReduce, glooAdaptorReduceScatter,

‎flagcx/adaptor/ccl/hccl_adaptor.cc‎

Lines changed: 21 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -141,6 +141,25 @@ flagcxResult_t hcclAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
141141
return flagcxNotSupported;
142142
}
143143

144+
// TODO: unsupported
145+
flagcxResult_t hcclAdaptorMemAlloc(void **ptr, size_t size) {
146+
return flagcxNotSupported;
147+
}
148+
149+
// TODO: unsupported
150+
flagcxResult_t hcclAdaptorMemFree(void *ptr) { return flagcxNotSupported; }
151+
152+
// TODO: unsupported
153+
flagcxResult_t hcclAdaptorCommRegister(flagcxInnerComm_t comm, void *buff,
154+
size_t size, void **handle) {
155+
return flagcxNotSupported;
156+
}
157+
158+
// TODO: unsupported
159+
flagcxResult_t hcclAdaptorCommDeregister(flagcxInnerComm_t comm, void *handle) {
160+
return flagcxNotSupported;
161+
}
162+
144163
flagcxResult_t hcclAdaptorReduce(const void *sendbuff, void *recvbuff,
145164
size_t count, flagcxDataType_t datatype,
146165
flagcxRedOp_t op, int root,
@@ -307,7 +326,8 @@ struct flagcxCCLAdaptor hcclAdaptor = {
307326
hcclAdaptorCommInitRank, hcclAdaptorCommFinalize, hcclAdaptorCommDestroy,
308327
hcclAdaptorCommAbort, hcclAdaptorCommResume, hcclAdaptorCommSuspend,
309328
hcclAdaptorCommCount, hcclAdaptorCommCuDevice, hcclAdaptorCommUserRank,
310-
hcclAdaptorCommGetAsyncError,
329+
hcclAdaptorCommGetAsyncError, hcclAdaptorMemAlloc, hcclAdaptorMemFree,
330+
hcclAdaptorCommRegister, hcclAdaptorCommDeregister,
311331
// Communication functions
312332
hcclAdaptorReduce, hcclAdaptorGather, hcclAdaptorScatter,
313333
hcclAdaptorBroadcast, hcclAdaptorAllReduce, hcclAdaptorReduceScatter,

‎flagcx/adaptor/ccl/ixnccl_adaptor.cc‎

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -72,6 +72,26 @@ flagcxResult_t ixncclAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
7272
(ncclResult_t *)&asyncError);
7373
}
7474

75+
// TODO: unsupported
76+
flagcxResult_t ixncclAdaptorMemAlloc(void **ptr, size_t size) {
77+
return flagcxNotSupported;
78+
}
79+
80+
// TODO: unsupported
81+
flagcxResult_t ixncclAdaptorMemFree(void *ptr) { return flagcxNotSupported; }
82+
83+
// TODO: unsupported
84+
flagcxResult_t ixncclAdaptorCommRegister(flagcxInnerComm_t comm, void *buff,
85+
size_t size, void **handle) {
86+
return flagcxNotSupported;
87+
}
88+
89+
// TODO: unsupported
90+
flagcxResult_t ixncclAdaptorCommDeregister(flagcxInnerComm_t comm,
91+
void *handle) {
92+
return flagcxNotSupported;
93+
}
94+
7595
flagcxResult_t ixncclAdaptorReduce(const void *sendbuff, void *recvbuff,
7696
size_t count, flagcxDataType_t datatype,
7797
flagcxRedOp_t op, int root,
@@ -262,6 +282,8 @@ struct flagcxCCLAdaptor ixncclAdaptor = {
262282
ixncclAdaptorCommDestroy, ixncclAdaptorCommAbort, ixncclAdaptorCommResume,
263283
ixncclAdaptorCommSuspend, ixncclAdaptorCommCount, ixncclAdaptorCommCuDevice,
264284
ixncclAdaptorCommUserRank, ixncclAdaptorCommGetAsyncError,
285+
ixncclAdaptorMemAlloc, ixncclAdaptorMemFree, ixncclAdaptorCommRegister,
286+
ixncclAdaptorCommDeregister,
265287
// Communication functions
266288
ixncclAdaptorReduce, ixncclAdaptorGather, ixncclAdaptorScatter,
267289
ixncclAdaptorBroadcast, ixncclAdaptorAllReduce, ixncclAdaptorReduceScatter,

‎flagcx/adaptor/ccl/mccl_adaptor.cc‎

Lines changed: 21 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -76,6 +76,25 @@ flagcxResult_t mcclAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
7676
(mcclResult_t *)&asyncError);
7777
}
7878

79+
// TODO: unsupported
80+
flagcxResult_t mcclAdaptorMemAlloc(void **ptr, size_t size) {
81+
return flagcxNotSupported;
82+
}
83+
84+
// TODO: unsupported
85+
flagcxResult_t mcclAdaptorMemFree(void *ptr) { return flagcxNotSupported; }
86+
87+
// TODO: unsupported
88+
flagcxResult_t mcclAdaptorCommRegister(flagcxInnerComm_t comm, void *buff,
89+
size_t size, void **handle) {
90+
return flagcxNotSupported;
91+
}
92+
93+
// TODO: unsupported
94+
flagcxResult_t mcclAdaptorCommDeregister(flagcxInnerComm_t comm, void *handle) {
95+
return flagcxNotSupported;
96+
}
97+
7998
flagcxResult_t mcclAdaptorReduce(const void *sendbuff, void *recvbuff,
8099
size_t count, flagcxDataType_t datatype,
81100
flagcxRedOp_t op, int root,
@@ -258,7 +277,8 @@ struct flagcxCCLAdaptor mcclAdaptor = {
258277
mcclAdaptorCommInitRank, mcclAdaptorCommFinalize, mcclAdaptorCommDestroy,
259278
mcclAdaptorCommAbort, mcclAdaptorCommResume, mcclAdaptorCommSuspend,
260279
mcclAdaptorCommCount, mcclAdaptorCommMcDevice, mcclAdaptorCommUserRank,
261-
mcclAdaptorCommGetAsyncError,
280+
mcclAdaptorCommGetAsyncError, mcclAdaptorMemAlloc, mcclAdaptorMemFree,
281+
mcclAdaptorCommRegister, mcclAdaptorCommDeregister,
262282
// Communication functions
263283
mcclAdaptorReduce, mcclAdaptorGather, mcclAdaptorScatter,
264284
mcclAdaptorBroadcast, mcclAdaptorAllReduce, mcclAdaptorReduceScatter,

‎flagcx/adaptor/ccl/mpi_adaptor.cc‎

Lines changed: 21 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -136,6 +136,25 @@ flagcxResult_t mpiAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
136136
return flagcxNotSupported;
137137
}
138138

139+
// TODO: unsupported
140+
flagcxResult_t mpiAdaptorMemAlloc(void **ptr, size_t size) {
141+
return flagcxNotSupported;
142+
}
143+
144+
// TODO: unsupported
145+
flagcxResult_t mpiAdaptorMemFree(void *ptr) { return flagcxNotSupported; }
146+
147+
// TODO: unsupported
148+
flagcxResult_t mpiAdaptorCommRegister(flagcxInnerComm_t comm, void *buff,
149+
size_t size, void **handle) {
150+
return flagcxNotSupported;
151+
}
152+
153+
// TODO: unsupported
154+
flagcxResult_t mpiAdaptorCommDeregister(flagcxInnerComm_t comm, void *handle) {
155+
return flagcxNotSupported;
156+
}
157+
139158
flagcxResult_t mpiAdaptorReduce(const void *sendbuff, void *recvbuff,
140159
size_t count, flagcxDataType_t datatype,
141160
flagcxRedOp_t op, int root,
@@ -321,7 +340,8 @@ struct flagcxCCLAdaptor mpiAdaptor = {
321340
mpiAdaptorCommInitRank, mpiAdaptorCommFinalize, mpiAdaptorCommDestroy,
322341
mpiAdaptorCommAbort, mpiAdaptorCommResume, mpiAdaptorCommSuspend,
323342
mpiAdaptorCommCount, mpiAdaptorCommCuDevice, mpiAdaptorCommUserRank,
324-
mpiAdaptorCommGetAsyncError,
343+
mpiAdaptorCommGetAsyncError, mpiAdaptorMemAlloc, mpiAdaptorMemFree,
344+
mpiAdaptorCommRegister, mpiAdaptorCommDeregister,
325345
// Communication functions
326346
mpiAdaptorReduce, mpiAdaptorGather, mpiAdaptorScatter, mpiAdaptorBroadcast,
327347
mpiAdaptorAllReduce, mpiAdaptorReduceScatter, mpiAdaptorAllGather,

‎flagcx/adaptor/ccl/musa_mccl_adaptor.cc‎

Lines changed: 21 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -72,6 +72,25 @@ flagcxResult_t mcclAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
7272
(mcclResult_t *)&asyncError);
7373
}
7474

75+
// TODO: unsupported
76+
flagcxResult_t mcclAdaptorMemAlloc(void **ptr, size_t size) {
77+
return flagcxNotSupported;
78+
}
79+
80+
// TODO: unsupported
81+
flagcxResult_t mcclAdaptorMemFree(void *ptr) { return flagcxNotSupported; }
82+
83+
// TODO: unsupported
84+
flagcxResult_t mcclAdaptorCommRegister(flagcxInnerComm_t comm, void *buff,
85+
size_t size, void **handle) {
86+
return flagcxNotSupported;
87+
}
88+
89+
// TODO: unsupported
90+
flagcxResult_t mcclAdaptorCommDeregister(flagcxInnerComm_t comm, void *handle) {
91+
return flagcxNotSupported;
92+
}
93+
7594
flagcxResult_t mcclAdaptorReduce(const void *sendbuff, void *recvbuff,
7695
size_t count, flagcxDataType_t datatype,
7796
flagcxRedOp_t op, int root,
@@ -253,7 +272,8 @@ struct flagcxCCLAdaptor musa_mcclAdaptor = {
253272
mcclAdaptorCommInitRank, mcclAdaptorCommFinalize, mcclAdaptorCommDestroy,
254273
mcclAdaptorCommAbort, mcclAdaptorCommResume, mcclAdaptorCommSuspend,
255274
mcclAdaptorCommCount, mcclAdaptorCommMuDevice, mcclAdaptorCommUserRank,
256-
mcclAdaptorCommGetAsyncError,
275+
mcclAdaptorCommGetAsyncError, mcclAdaptorMemAlloc, mcclAdaptorMemFree,
276+
mcclAdaptorCommRegister, mcclAdaptorCommDeregister,
257277
// Communication functions
258278
mcclAdaptorReduce, mcclAdaptorGather, mcclAdaptorScatter,
259279
mcclAdaptorBroadcast, mcclAdaptorAllReduce, mcclAdaptorReduceScatter,

‎flagcx/adaptor/ccl/nccl_adaptor.cc‎

Lines changed: 20 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -71,6 +71,24 @@ flagcxResult_t ncclAdaptorCommGetAsyncError(flagcxInnerComm_t comm,
7171
(ncclResult_t *)&asyncError);
7272
}
7373

74+
flagcxResult_t ncclAdaptorMemAlloc(void **ptr, size_t size) {
75+
return (flagcxResult_t)ncclMemAlloc(ptr, size);
76+
}
77+
78+
flagcxResult_t ncclAdaptorMemFree(void *ptr) {
79+
return (flagcxResult_t)ncclMemFree(ptr);
80+
}
81+
82+
flagcxResult_t ncclAdaptorCommRegister(const flagcxInnerComm_t comm, void *buff,
83+
size_t size, void **handle) {
84+
return (flagcxResult_t)ncclCommRegister(comm->base, buff, size, handle);
85+
}
86+
87+
flagcxResult_t ncclAdaptorCommDeregister(const flagcxInnerComm_t comm,
88+
void *handle) {
89+
return (flagcxResult_t)ncclCommDeregister(comm->base, handle);
90+
}
91+
7492
flagcxResult_t ncclAdaptorReduce(const void *sendbuff, void *recvbuff,
7593
size_t count, flagcxDataType_t datatype,
7694
flagcxRedOp_t op, int root,
@@ -252,7 +270,8 @@ struct flagcxCCLAdaptor ncclAdaptor = {
252270
ncclAdaptorCommInitRank, ncclAdaptorCommFinalize, ncclAdaptorCommDestroy,
253271
ncclAdaptorCommAbort, ncclAdaptorCommResume, ncclAdaptorCommSuspend,
254272
ncclAdaptorCommCount, ncclAdaptorCommCuDevice, ncclAdaptorCommUserRank,
255-
ncclAdaptorCommGetAsyncError,
273+
ncclAdaptorCommGetAsyncError, ncclAdaptorMemAlloc, ncclAdaptorMemFree,
274+
ncclAdaptorCommRegister, ncclAdaptorCommDeregister,
256275
// Communication functions
257276
ncclAdaptorReduce, ncclAdaptorGather, ncclAdaptorScatter,
258277
ncclAdaptorBroadcast, ncclAdaptorAllReduce, ncclAdaptorReduceScatter,

0 commit comments

Comments
 (0)