3. 编程模型

3.1. 通讯器管理

3.1.1. mcclGetVersion

mcclResult_t mcclGetVersion(int* version);

获取当前MCCL版本号。

3.1.2. mcclGetUniqueId

mcclResult_t mcclGetUniqueId(mcclUniqueId* uniqueId);

获取ID号，用于在多进程场景中初始化的ID，由任意一个进程获取，并通过MPI或socket发送给其它进程用于初始化。多节点环境同样适用。

3.1.3. mcclCommInitRank

mcclResult_t mcclCommInitRank(mcclComm_t* comm, int nranks, mcclUniqueId commId, int rank);

创建通讯器， nranks 为通讯器总数， commId 为使用 mcclGetUniqueId 获取到的ID， rank 为此通讯器的编号，其最小值为 0，最大值为 nranks-1。此通讯器所使用的设备为当前线程正在使用的设备。每个通讯器不能使用相同的设备。

3.1.4. mcclCommInitAll

mcclResult_t mcclCommInitAll(mcclComm_t* comm, int ndev, const int* devlist);

用于在单进程场景中一次性创建所有通讯器。 devlist 包含了每个rank使用的设备。

3.1.5. mcclCommDestroy

mcclResult_t mcclCommDestroy(mcclComm_t comm);

销毁通讯器，结束集合通讯，用于正常通讯流程的退出。

3.1.6. mcclCommAbort

mcclResult_t mcclCommAbort(mcclComm_t comm);

中止通讯器，结束集合通讯，常用于从致命异常（如网络故障、节点故障或进程故障）中恢复通讯。当异常出现时，调用此API，然后创建新的通信器重新通讯。

3.1.7. mcclGetErrorString

const char* mcclGetErrorString(mcclResult_t result);

获取返回值的字符串。

3.1.8. mcclCommGetAsyncError

mcclResult_t mcclCommGetAsyncError(mcclComm_t comm, mcclResult_t* asyncError);

获取异步操作的返回值。

3.1.9. mcclCommCount

mcclResult_t mcclCommCount(const mcclComm_t comm, int* count);

获取通讯器数量。

3.1.10. mcclCommMcDevice

mcclResult_t mcclCommMcDevice(const mcclComm_t comm, int* device);

获取当前通讯器所使用的 device。

3.1.11. mcclCommUserRank

mcclResult_t mcclCommUserRank(const mcclComm_t comm, int* rank);

获取当前通讯器的 rank 编号。

3.1.12. 示例

3.1.12.1. 单进程

int nranks = 4;
mcclComm_t comms[nranks];
int devs[nranks] = {0, 1, 2, 3};
mcclCommInitAll(comms, nranks, devs);
//...
//OPs...
//...
// finalizing MCCL
for (int i = 0; i < nranks; ++i)
   mcclCommDestroy(comms[i]);

3.1.12.2. 多进程

mcclUniqueId id;
mcclComm_t comm;

//initializing MPI
MPICHECK(MPI_Init(&argc, &argv));
MPICHECK(MPI_Comm_rank(MPI_COMM_WORLD, &myRank));
MPICHECK(MPI_Comm_size(MPI_COMM_WORLD, &nRanks));

//get MCCL unique ID at rank 0 and broadcast it to all others
if (myRank == 0) mcclGetUniqueId(&id);
MPICHECK(MPI_Bcast((void *)&id, sizeof(id), MPI_BYTE, 0, MPI_COMM_WORLD));

mcSetDevice(myRank);
mcclCommInitRank(&comm, nRanks, id, myRank);

//Other Ops.

//finalizing MCCL
mcclCommDestroy(comm);

3.2. 集合通讯

所有rank都必须参与集合通讯，即每一种操作必须被所有rank调用，否则其它rank将等待。

3.2.1. mcclAllReduce

mcclResult_t mcclAllReduce(const void* sendbuff, void* recvbuff, size_t count, mcclDataType_t datatype, mcclRedOp_t op, mcclComm_t comm, mcStream_t stream);

在所有rank上都执行归约计算，最终所有rank将得到相同的结果，并保存在 recvbuff 中。 sendbuff 为输入源， count 为数据的个数， datatype 为数据类型， op 为操作类型， comm 为通信器对象， stream 为MXMACA的stream。

3.2.2. mcclBroadcast

mcclResult_t mcclBroadcast(const void* sendbuff, void* recvbuff, size_t count, mcclDataType_t datatype, int root, mcclComm_t comm, mcStream_t stream);

mcclResult_t mcclBcast(void* buff, size_t count, mcclDataType_t datatype, int root, mcclComm_t comm, mcStream_t stream);

mcclBroadcast 可将 sendbuff 中的数据复制到所有rank的 recvbuff 中（包括发送者的 recvbuff ）。 count 为数据的个数， datatype 为数据类型， root 为发送者的rank编号， comm 为通信器对象， stream 为MXMACA的stream，对于接收者， sendbuff 可为空。

mcclBcast 与 mcclBroadcast 相同，若发送者的 sendbuff 与 recvbuff 相同，则可以用 mcclBcast 代替 mcclBroadcast 。

3.2.3. mcclReduce

mcclResult_t mcclReduce(const void* sendbuff, void* recvbuff, size_t count, mcclDataType_t datatype, mcclRedOp_t op, int root, mcclComm_t comm, mcStream_t stream);

在rank为 root 的通讯器上执行归约计算，结果保存在 root 的 recvbuff 中。对于其它rank， recvbuff 可为空。对于发送方， root 为发送的目标rank。 root 为执行归约的通讯器rank号， comm 为通信器对象， stream 为MXMACA的stream。

3.2.4. mcclAllGather

mcclResult_t mcclAllGather(const void* sendbuff, void* recvbuff, size_t sendcount, mcclDataType_t datatype, mcclComm_t comm, mcStream_t stream);

将所有rank的 sendbuff 中的数据按rank顺序排列保存到每一个rank的 recvbuff 中，最终每个rank得到相同的结果。 sendcount 为 sendbuff 的数据个数， recvbuff 的数据个数将会是 sendcount*N （N 为所有rank的个数）。 datatype 为数据类型， comm 为通信器对象， stream 为MXMACA的stream。调用者必须确保 recvbuff 有足够的空间。

3.2.5. mcclReduceScatter

mcclResult_t mcclReduceScatter(const void* sendbuff, void* recvbuff, size_t recvcount, mcclDataType_t datatype, mcclRedOp_t op, mcclComm_t comm, mcStream_t stream);

对所有rank的 sendbuff 进行归约，然后将结果分成 N 份，并按rank顺序分发给所有通讯器的 recvbuff 中。 recvcount 为接收数据的个数，发送数据的个数是 recvcount*N 。 op 为操作类型， comm 为通信器对象， stream 为MXMACA的stream。

3.2.6. mcclAllToAll

mcclResult_t mcclAllToAll(const void* sendbuff, void* recvbuff, size_t count, mcclDataType_t datatype, mcclComm_t comm, mcStream_t stream);

在所有rank上执行全交换操作，每个rank将 sendbuff 中的数据发送给所有其他rank，并从其他rank接收数据到 recvbuff 中。每个rank发送和接收的数据量相同，均为 count 个 datatype 类型的数据。最终，每个rank的 recvbuff 中将包含来自所有rank的数据，且数据顺序与rank编号一致。 sendbuff 为发送数据的缓冲区， recvbuff 为接收数据的缓冲区， count 为每个rank发送和接收的数据个数。 datatype 为数据类型， comm 为通信器对象， stream 为MXMACA的stream。调用者必须确保 recvbuff 有足够的空间来存储所有rank发送的数据。

3.2.7. mcclAllToAllv

mcclResult_t mcclAllToAllv(const void* sendbuff, const size_t sendcounts[], const size_t sdispls[], void* recvbuff, const size_t recvcounts[], const size_t rdispls[], mcclDataType_t datatype, mcclComm_t comm, mcStream_t stream);

在所有rank上执行非均匀全交换操作，每个rank可以发送不同数量的数据给其他rank，并从其他rank接收不同数量的数据。 sendbuff 为发送数据的缓冲区， sendcounts 为数组， sendcounts[i] 表示发送给rank i 的数据个数。 sdispls 为数组， sdispls[i] 表示发送给rank i 的数据在sendbuff中的起始偏移量。 recvbuff 为接收数据的缓冲区， recvcounts 为数组， recvcounts[i] 表示从rank i 接收的数据个数。 rdispls 为数组， rdispls[i] 表示从rank i 接收的数据在 recvbuff 中的起始偏移量。 datatype 为数据类型， comm 为通信器对象， stream 为MXMACA的stream。调用者必须确保 sendbuff 和 recvbuff 有足够的空间，且 sendcounts、 sdispls 、 recvcounts 和 rdispls 的数组长度必须与通信器中的rank数量一致。

3.2.8. mcclAllToAlld

mcclResult_t mcclAllToAlld(const void* sendbuff[], const size_t sendcounts[], void* recvbuff[], const size_t recvcounts[], mcclDataType_t datatype, mcclComm_t comm, mcStream_t stream);

在所有rank上执行分布式全交换操作，每个rank发送多个数据块给其他rank，并从其他rank接收多个数据块。 sendbuff 和 recvbuff 是数组的数组，分别表示每个rank发送和接收的数据块。 sendcounts 和 recvcounts 分别指定每个rank发送和接收的数据块的大小。 sendbuff 为数组， sendbuff[i] 指向发送给rank i 的数据缓冲区， sendcounts 为数组， sendcounts[i] 表示发送给rank i 的数据个数。 recvbuff 为数组， recvbuff[i] 指向接收来自rank i 的数据缓冲区， recvcounts 为数组， recvcounts[i] 表示从rank i 接收的数据个数。 datatype 为数据类型， comm 为通信器对象， stream 为MXMACA的stream。调用者必须确保 sendbuff 和 recvbuff 中的每个缓冲区有足够的空间，且 sendbuff、 recvbuff 、 sendcounts 和 recvcounts 的数组长度必须与通信器中的rank数量一致。

3.2.9. 示例

mcclGroupStart();
for (int i = 0; i < nranks; ++i) {
   mcclAllReduce((const void*)sendbuff[i], (void*)recvbuff[i], size, mcclFloat, mcclSum, comms[i], s[i]);
}
mcclGroupEnd();

mcclGroupStart();
for (int i = 0; i < nranks; ++i) {
   mcclBroadcast((const void*)sendbuff[i], (void*)recvbuff[i], size, mcclFloat, 0, comms[i], s[i]);
}
mcclGroupEnd();

mcclGroupStart();
for (int i = 0; i < nranks; ++i) {
   mcclReduce((const void*)sendbuff[i], (void*)recvbuff[i], size, mcclFloat, mcclSum, 0, comms[i], s[i]);
}
mcclGroupEnd();

mcclGroupStart();
for (int i = 0; i < nranks; ++i) {
   mcclAllGather((const void*)sendbuff[i], (void*)recvbuff[i], size, mcclFloat, comms[i], s[i]);
}
mcclGroupEnd();

mcclGroupStart();
for (int i = 0; i < nranks; ++i) {
   mcclAllReduceScatter((const void*)sendbuff[i], (void*)recvbuff[i], size, mcclFloat, mcclSum, comms[i], s[i]);
}
mcclGroupEnd();

mcclGroupStart();
for (int i = 0; i < nranks; ++i) {
   mcclAllToAll((const void*)sendbuff[i], (void*)recvbuff[i], size, mcclFloat, comms[i], s[i]);
}
mcclGroupEnd();

mcclGroupStart();
for (int i = 0; i < nranks; ++i) {
   mcclToAllAllv((const void*)sendbuff[i], sendcounts, sdispls, (void*)recvbuff[i], recvcounts, rdispls, mcclFloat, comms[i], s[i]);
}
mcclGroupEnd();

mcclGroupStart();
for (int i = 0; i < nranks; ++i) {
   mcclAllToAlld((const void**)sendbuff[i], sendcounts, (void**)recvbuff[i], recvcounts, mcclFloat, comms[i], s[i]);
}
mcclGroupEnd();

3.3. 组调用

组调用可以把一系列MCCL操作合并在一起放入后台执行，不同stream的操作可互不影响地执行。这样可以实现MCCL操作在不同stream并发执行。

3.3.1. mcclGroupStart

mcclResult_t mcclGroupStart();

开始一个组编程。在此之后的所有MCCL函数都不会立即执行，而是放入一个执行队列中等待执行。

3.3.2. mcclGroupEnd

mcclResult_t mcclGroupEnd();

结束组编程。立即启动队列中的所有操作，并根据不同stream放入不同队列执行。在所有操作完成之前，阻塞当前线程。

3.3.3. 示例

mcclGroupStart();
mcclSend(sendbuff[0], size, mcclFloat, 1, comms[0], s[0]);
mcclRecv(recvbuff[1], size, mcclFloat, 0, comms[1], s[1]);
mcclGroupEnd();

3.4. 点对点通讯

点对点的通讯，即把数据从一个rank复制给另一个rank，也就是将显存中的内容，由一个显卡复制到另一个显卡中。

3.4.1. mcclSend

mcclResult_t mcclSend(const void* sendbuff, size_t count, mcclDataType_t datatype, int peer, mcclComm_t comm, mcStream_t stream);

将 sendbuff 中的内存发送给接收端 peer 。 count 为数据的个数， datatype 为数据类型。

comm 为通信器对象， stream 为MXMACA的stream。

3.4.2. mcclRecv

mcclResult_t mcclRecv(void* recvbuff, size_t count, mcclDataType_t datatype, int peer, mcclComm_t comm, mcStream _t stream);

从 peer 端接收数据到 recvbuff 中。 count 为数据的个数， datatype 为数据类型。

comm 为通信器对象， stream 为MXMACA的stream。

3.4.3. 示例

mcclSend(sendbuff[0], size, mcclFloat, 1, comms[0], s[0]);
mcclRecv(recvbuff[1], size, mcclFloat, 0, comms[1], s[1]);

3.5. 数据类型

MCCL的数据类型用于指定各操作原语所使用的数据类型，对应MXMACA的数据类型。对应关系如下：

mcclInt8：char
mcclChar：char
mcclUint8：unsigned char
mcclInt32：int
mcclInt：int
mcclUint32：unsigned int
mcclInt64：long
mcclUint64：unsigned long
mcclFloat16：half
mcclHalf：half
mcclFloat32：float
mcclFloat：float
mcclFloat64：double
mcclDouble：double
mcclBfloat16：bfloat16

3.6. 返回值类型

mcclSuccess：成功
mcclUnhandledMacaError：调用MXMACA API时返回错误
mcclSystemError：调用系统函数错误
mcclInternalError：MCCL内部错误，可能是内存错误
mcclInvalidArgument：参数错误
mcclInvalidUsage：使用错误，可能是用户编程导致
mcclRemoteError：多节点场景时，网络连接错误

3.7. 归约操作类型

mcclSum：求和
mcclProd：求乘积
mcclMax：求最大值
mcclMin：求最小值
mcclAvg：求平均值