Bertan Dogancay
374 baris
15 KiB
C++
374 baris
15 KiB
C++
/*************************************************************************
|
|
* Copyright (c) 2015-2023, NVIDIA CORPORATION. All rights reserved.
|
|
*
|
|
* See LICENSE.txt for license information
|
|
************************************************************************/
|
|
|
|
#include "argcheck.h" // Need some checks here since we access comm
|
|
#include "collectives.h"
|
|
#include "enqueue.h"
|
|
#include "graph/topo.h"
|
|
#include "nccl.h"
|
|
|
|
#include "msccl/msccl_lifecycle.h"
|
|
|
|
NCCL_API(ncclResult_t, ncclAllGather, const void* sendbuff, void* recvbuff, size_t sendcount,
|
|
ncclDataType_t datatype, ncclComm_t comm, cudaStream_t stream);
|
|
ncclResult_t ncclAllGather(const void* sendbuff, void* recvbuff, size_t sendcount,
|
|
ncclDataType_t datatype, ncclComm_t comm, cudaStream_t stream) {
|
|
// Just pass the size of one message and not the total bytes sent/received.
|
|
constexpr nvtxPayloadSchemaEntry_t AllGatherSchema[] = {
|
|
{0, NVTX_PAYLOAD_ENTRY_TYPE_SIZE, "Message size [bytes]"}
|
|
};
|
|
size_t msgsize = sendcount * ncclTypeSize(datatype);
|
|
NVTX3_FUNC_WITH_PARAMS(AllGather, AllGatherSchema, msgsize)
|
|
|
|
if (mscclAvailable() && !mscclIsCaller()) {
|
|
return mscclEnqueueCheck(
|
|
sendbuff, nullptr, nullptr, recvbuff, nullptr, nullptr,
|
|
sendcount, datatype, 0, 0, ncclSum, mscclFuncAllGather, comm, stream);
|
|
}
|
|
|
|
struct ncclInfo info = { ncclFuncAllGather, "AllGather",
|
|
sendbuff, recvbuff, sendcount, datatype, ncclSum, 0, comm, stream, /* Args */
|
|
ALLGATHER_CHUNKSTEPS, ALLGATHER_SLICESTEPS };
|
|
return ncclEnqueueCheck(&info);
|
|
}
|
|
|
|
NCCL_API(ncclResult_t, ncclAllReduce, const void* sendbuff, void* recvbuff, size_t count,
|
|
ncclDataType_t datatype, ncclRedOp_t op, ncclComm* comm, cudaStream_t stream);
|
|
ncclResult_t ncclAllReduce(const void* sendbuff, void* recvbuff, size_t count,
|
|
ncclDataType_t datatype, ncclRedOp_t op, ncclComm* comm, cudaStream_t stream) {
|
|
struct NvtxParamsAllReduce {
|
|
size_t bytes;
|
|
ncclRedOp_t op;
|
|
};
|
|
// Just pass the size of one message and not the total bytes sent/received.
|
|
static constexpr nvtxPayloadSchemaEntry_t AllReduceSchema[] = {
|
|
{0, NVTX_PAYLOAD_ENTRY_TYPE_SIZE, "Message size [bytes]"},
|
|
{0, NVTX_PAYLOAD_ENTRY_NCCL_REDOP, "Reduction operation", nullptr, 0,
|
|
offsetof(NvtxParamsAllReduce, op)}
|
|
};
|
|
NvtxParamsAllReduce payload{count * ncclTypeSize(datatype), op};
|
|
NVTX3_FUNC_WITH_PARAMS(AllReduce, AllReduceSchema, payload)
|
|
|
|
if (mscclAvailable() && !mscclIsCaller()) {
|
|
return mscclEnqueueCheck(
|
|
sendbuff, nullptr, nullptr, recvbuff, nullptr, nullptr,
|
|
count, datatype, 0, 0, op, mscclFuncAllReduce, comm, stream);
|
|
}
|
|
|
|
struct ncclInfo info = { ncclFuncAllReduce, "AllReduce",
|
|
sendbuff, recvbuff, count, datatype, op, 0, comm, stream, /* Args */
|
|
ALLREDUCE_CHUNKSTEPS, ALLREDUCE_SLICESTEPS };
|
|
return ncclEnqueueCheck(&info);
|
|
}
|
|
|
|
NCCL_API(ncclResult_t, ncclAllToAll, const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype,
|
|
ncclComm_t comm, hipStream_t stream);
|
|
ncclResult_t ncclAllToAll(const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype,
|
|
ncclComm_t comm, hipStream_t stream) {
|
|
// Just pass the size of one message and not the total bytes sent/received.
|
|
constexpr nvtxPayloadSchemaEntry_t AllToAllSchema[] = {
|
|
{0, NVTX_PAYLOAD_ENTRY_TYPE_SIZE, "Message size [bytes]"}
|
|
};
|
|
size_t msgsize = count * ncclTypeSize(datatype);
|
|
NVTX3_FUNC_WITH_PARAMS(AllToAll, AllToAllSchema, msgsize)
|
|
|
|
if (mscclAvailable() && !mscclIsCaller()) {
|
|
return mscclEnqueueCheck(
|
|
sendbuff, nullptr, nullptr, recvbuff, nullptr, nullptr,
|
|
count, datatype, 0, 0, ncclSum, mscclFuncAllToAll, comm, stream);
|
|
}
|
|
|
|
size_t rankOffset = count * ncclTypeSize(datatype);
|
|
size_t rankAlign = rankOffset & ((~rankOffset) + 1);
|
|
// Determine Pivot A2A support now that we know number of channels
|
|
if (comm->topo->pivotA2AEnabled && comm->nChannels >= comm->topo->pivotA2ANumBiRings * 2 &&
|
|
rankOffset >= 744 * 1024 && rankAlign != 4) {
|
|
struct ncclInfo info = { ncclFuncAllToAllPivot, "AllToAllPivot",
|
|
sendbuff, recvbuff, count, datatype, ncclSum, 0, comm, stream, /* Args */
|
|
ALLTOALL_PIVOT_CHUNKSTEPS, ALLTOALL_PIVOT_SLICESTEPS };
|
|
return ncclEnqueueCheck(&info);
|
|
} else {
|
|
int nRanks;
|
|
NCCLCHECK(ncclCommCount(comm, &nRanks));
|
|
if (count == 0) return ncclSuccess;
|
|
NCCLCHECK(ncclGroupStart());
|
|
for (int r=0; r<nRanks; r++) {
|
|
NCCLCHECK(ncclSend(((char*)sendbuff)+r*rankOffset, count, datatype, r, comm, stream));
|
|
NCCLCHECK(ncclRecv(((char*)recvbuff)+r*rankOffset, count, datatype, r, comm, stream));
|
|
}
|
|
NCCLCHECK(ncclGroupEnd());
|
|
return ncclSuccess;
|
|
}
|
|
}
|
|
|
|
NCCL_API(ncclResult_t, ncclAllToAllv, const void *sendbuff, const size_t sendcounts[], const size_t sdispls[],
|
|
void *recvbuff, const size_t recvcounts[], const size_t rdispls[],
|
|
ncclDataType_t datatype, ncclComm_t comm, hipStream_t stream);
|
|
ncclResult_t ncclAllToAllv(const void *sendbuff, const size_t sendcounts[], const size_t sdispls[],
|
|
void *recvbuff, const size_t recvcounts[], const size_t rdispls[],
|
|
ncclDataType_t datatype, ncclComm_t comm, hipStream_t stream) {
|
|
struct NvtxParamsAllToAllv {
|
|
size_t sendbytes;
|
|
size_t recvbytes;
|
|
};
|
|
// Just pass the size of one send/recv messages and not the total bytes sent/received.
|
|
constexpr nvtxPayloadSchemaEntry_t AllToAllvSchema[] = {
|
|
{0, NVTX_PAYLOAD_ENTRY_TYPE_SIZE, "Message size [bytes] (Send)"},
|
|
{0, NVTX_PAYLOAD_ENTRY_TYPE_SIZE, "Message size [bytes] (Recv)"}
|
|
};
|
|
NvtxParamsAllToAllv payload{sendcounts[comm->rank] * ncclTypeSize(datatype), recvcounts[comm->rank] * ncclTypeSize(datatype)};
|
|
NVTX3_FUNC_WITH_PARAMS(AllToAllv, AllToAllvSchema, payload)
|
|
|
|
if (mscclAvailable() && !mscclIsCaller()) {
|
|
return mscclEnqueueCheck(
|
|
sendbuff, sendcounts, sdispls, recvbuff, recvcounts, rdispls,
|
|
0, datatype, 0, 0, ncclSum, mscclFuncAllToAllv, comm, stream);
|
|
}
|
|
|
|
int nRanks;
|
|
NCCLCHECK(ncclCommCount(comm, &nRanks));
|
|
NCCLCHECK(ncclGroupStart());
|
|
for (int r=0; r<nRanks; r++) {
|
|
if (sendcounts[r]) NCCLCHECK(ncclSend(
|
|
((char*)sendbuff) + sdispls[r]*ncclTypeSize(datatype),
|
|
sendcounts[r],
|
|
datatype,
|
|
r,
|
|
comm,
|
|
stream));
|
|
if (recvcounts[r]) NCCLCHECK(ncclRecv(
|
|
((char*)recvbuff) + rdispls[r]*ncclTypeSize(datatype),
|
|
recvcounts[r],
|
|
datatype,
|
|
r,
|
|
comm,
|
|
stream));
|
|
}
|
|
NCCLCHECK(ncclGroupEnd());
|
|
return ncclSuccess;
|
|
}
|
|
|
|
NCCL_API(ncclResult_t, ncclBroadcast, const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype, int root,
|
|
ncclComm_t comm, cudaStream_t stream);
|
|
ncclResult_t ncclBroadcast(const void* sendbuff, void* recvbuff, size_t count, ncclDataType_t datatype, int root,
|
|
ncclComm_t comm, cudaStream_t stream) {
|
|
struct NvtxParamsBroadcast {
|
|
size_t bytes;
|
|
int root;
|
|
};
|
|
constexpr nvtxPayloadSchemaEntry_t BroadcastSchema[] = {
|
|
{0, NVTX_PAYLOAD_ENTRY_TYPE_SIZE, "Bytes"},
|
|
{0, NVTX_PAYLOAD_ENTRY_TYPE_INT, "Root", nullptr, 0, offsetof(NvtxParamsBroadcast, root)}
|
|
};
|
|
NvtxParamsBroadcast payload{count * ncclTypeSize(datatype), root};
|
|
NVTX3_FUNC_WITH_PARAMS(Broadcast, BroadcastSchema, payload)
|
|
|
|
if (mscclAvailable() && !mscclIsCaller()) {
|
|
return mscclEnqueueCheck(
|
|
sendbuff, nullptr, nullptr, recvbuff, nullptr, nullptr,
|
|
count, datatype, root, 0, ncclSum, mscclFuncBroadcast, comm, stream);
|
|
}
|
|
|
|
struct ncclInfo info = { ncclFuncBroadcast, "Broadcast",
|
|
sendbuff, recvbuff, count, datatype, ncclSum, root, comm, stream, /* Args */
|
|
BROADCAST_CHUNKSTEPS, BROADCAST_SLICESTEPS };
|
|
return ncclEnqueueCheck(&info);
|
|
}
|
|
/* Deprecated original "in place" function, similar to MPI */
|
|
NCCL_API(ncclResult_t, ncclBcast, void* buff, size_t count, ncclDataType_t datatype, int root,
|
|
ncclComm_t comm, cudaStream_t stream);
|
|
ncclResult_t ncclBcast(void* buff, size_t count, ncclDataType_t datatype, int root,
|
|
ncclComm_t comm, cudaStream_t stream) {
|
|
return ncclBroadcast(buff, buff, count, datatype, root, comm, stream);
|
|
}
|
|
|
|
NCCL_API(ncclResult_t, ncclGather, const void* sendbuff, void* recvbuff, size_t sendcount,
|
|
ncclDataType_t datatype, int root, ncclComm_t comm, hipStream_t stream);
|
|
ncclResult_t ncclGather(const void* sendbuff, void* recvbuff, size_t sendcount,
|
|
ncclDataType_t datatype, int root, ncclComm_t comm, hipStream_t stream) {
|
|
struct NvtxParamsGather {
|
|
size_t bytes;
|
|
int root;
|
|
};
|
|
constexpr nvtxPayloadSchemaEntry_t GatherSchema[] = {
|
|
{0, NVTX_PAYLOAD_ENTRY_TYPE_SIZE, "Bytes"},
|
|
{0, NVTX_PAYLOAD_ENTRY_TYPE_INT, "Root", nullptr, 0, offsetof(NvtxParamsGather, root)}
|
|
};
|
|
NvtxParamsGather payload{sendcount * ncclTypeSize(datatype), root};
|
|
NVTX3_FUNC_WITH_PARAMS(Gather, GatherSchema, payload)
|
|
|
|
if (mscclAvailable() && !mscclIsCaller()) {
|
|
return mscclEnqueueCheck(
|
|
sendbuff, nullptr, nullptr, recvbuff, nullptr, nullptr,
|
|
sendcount, datatype, root, 0, ncclSum, mscclFuncGather, comm, stream);
|
|
}
|
|
|
|
int nRanks;
|
|
NCCLCHECK(ncclCommCount(comm, &nRanks));
|
|
size_t rankOffset = sendcount * ncclTypeSize(datatype);
|
|
if (sendcount == 0) return ncclSuccess;
|
|
int rank;
|
|
NCCLCHECK(ncclCommUserRank(comm, &rank));
|
|
NCCLCHECK(ncclGroupStart());
|
|
if (rank == root) {
|
|
for (int r=0; r<nRanks; r++)
|
|
NCCLCHECK(ncclRecv(((char*)recvbuff)+r*rankOffset, sendcount, datatype, r, comm, stream));
|
|
}
|
|
NCCLCHECK(ncclSend(sendbuff, sendcount, datatype, root, comm, stream));
|
|
NCCLCHECK(ncclGroupEnd());
|
|
return ncclSuccess;
|
|
}
|
|
|
|
NCCL_API(ncclResult_t, ncclReduce, const void* sendbuff, void* recvbuff, size_t count,
|
|
ncclDataType_t datatype, ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream);
|
|
ncclResult_t ncclReduce(const void* sendbuff, void* recvbuff, size_t count,
|
|
ncclDataType_t datatype, ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream) {
|
|
struct NvtxParamsReduce {
|
|
size_t bytes;
|
|
int root;
|
|
ncclRedOp_t op;
|
|
};
|
|
constexpr nvtxPayloadSchemaEntry_t ReduceSchema[] = {
|
|
{0, NVTX_PAYLOAD_ENTRY_TYPE_SIZE, "Message size [bytes]"},
|
|
{0, NVTX_PAYLOAD_ENTRY_TYPE_INT, "Root", nullptr, 0, offsetof(NvtxParamsReduce, root)},
|
|
{0, NVTX_PAYLOAD_ENTRY_NCCL_REDOP, "Reduction operation", nullptr, 0,
|
|
offsetof(NvtxParamsReduce, op)}
|
|
};
|
|
NvtxParamsReduce payload{count * ncclTypeSize(datatype), root, op};
|
|
NVTX3_FUNC_WITH_PARAMS(Reduce, ReduceSchema, payload)
|
|
|
|
if (mscclAvailable() && !mscclIsCaller()) {
|
|
return mscclEnqueueCheck(
|
|
sendbuff, nullptr, nullptr, recvbuff, nullptr, nullptr,
|
|
count, datatype, root, 0, op, mscclFuncReduce, comm, stream);
|
|
}
|
|
|
|
struct ncclInfo info = { ncclFuncReduce, "Reduce",
|
|
sendbuff, recvbuff, count, datatype, op, root, comm, stream, /* Args */
|
|
REDUCE_CHUNKSTEPS, REDUCE_SLICESTEPS };
|
|
return ncclEnqueueCheck(&info);
|
|
}
|
|
|
|
NCCL_API(ncclResult_t, ncclReduceScatter, const void* sendbuff, void* recvbuff, size_t recvcount,
|
|
ncclDataType_t datatype, ncclRedOp_t op, ncclComm* comm, cudaStream_t stream);
|
|
ncclResult_t ncclReduceScatter(const void* sendbuff, void* recvbuff, size_t recvcount,
|
|
ncclDataType_t datatype, ncclRedOp_t op, ncclComm* comm, cudaStream_t stream) {
|
|
struct NvtxParamsReduceScatter {
|
|
size_t bytes;
|
|
ncclRedOp_t op;
|
|
};
|
|
constexpr nvtxPayloadSchemaEntry_t ReduceScatterSchema[] = {
|
|
{0, NVTX_PAYLOAD_ENTRY_TYPE_SIZE, "Message size [bytes]"},
|
|
{0, NVTX_PAYLOAD_ENTRY_NCCL_REDOP, "Reduction operation", nullptr, 0,
|
|
offsetof(NvtxParamsReduceScatter, op)}
|
|
};
|
|
NvtxParamsReduceScatter payload{recvcount * ncclTypeSize(datatype), op};
|
|
NVTX3_FUNC_WITH_PARAMS(ReduceScatter, ReduceScatterSchema, payload)
|
|
|
|
if (mscclAvailable() && !mscclIsCaller()) {
|
|
return mscclEnqueueCheck(
|
|
sendbuff, nullptr, nullptr, recvbuff, nullptr, nullptr,
|
|
recvcount, datatype, 0, 0, op, mscclFuncReduceScatter, comm, stream);
|
|
}
|
|
|
|
struct ncclInfo info = { ncclFuncReduceScatter, "ReduceScatter",
|
|
sendbuff, recvbuff, recvcount, datatype, op, 0, comm, stream, /* Args */
|
|
REDUCESCATTER_CHUNKSTEPS, REDUCESCATTER_SLICESTEPS };
|
|
return ncclEnqueueCheck(&info);
|
|
}
|
|
|
|
NCCL_API(ncclResult_t, ncclScatter, const void* sendbuff, void* recvbuff, size_t recvcount, ncclDataType_t datatype, int root,
|
|
ncclComm_t comm, hipStream_t stream);
|
|
ncclResult_t ncclScatter(const void* sendbuff, void* recvbuff, size_t recvcount, ncclDataType_t datatype, int root,
|
|
ncclComm_t comm, hipStream_t stream) {
|
|
struct NvtxParamsScatter {
|
|
size_t bytes;
|
|
int root;
|
|
};
|
|
constexpr nvtxPayloadSchemaEntry_t ScatterSchema[] = {
|
|
{0, NVTX_PAYLOAD_ENTRY_TYPE_SIZE, "Bytes"},
|
|
{0, NVTX_PAYLOAD_ENTRY_TYPE_INT, "Root", nullptr, 0, offsetof(NvtxParamsScatter, root)}
|
|
};
|
|
NvtxParamsScatter payload{recvcount * ncclTypeSize(datatype), root};
|
|
NVTX3_FUNC_WITH_PARAMS(Scatter, ScatterSchema, payload)
|
|
|
|
if (mscclAvailable() && !mscclIsCaller()) {
|
|
return mscclEnqueueCheck(
|
|
sendbuff, nullptr, nullptr, recvbuff, nullptr, nullptr,
|
|
recvcount, datatype, root, 0, ncclSum, mscclFuncScatter, comm, stream);
|
|
}
|
|
|
|
int nRanks;
|
|
NCCLCHECK(ncclCommCount(comm, &nRanks));
|
|
size_t rankOffset = recvcount * ncclTypeSize(datatype);
|
|
if (recvcount == 0) return ncclSuccess;
|
|
int rank;
|
|
NCCLCHECK(ncclCommUserRank(comm, &rank));
|
|
NCCLCHECK(ncclGroupStart());
|
|
if (rank == root) {
|
|
for (int r=0; r<nRanks; r++)
|
|
NCCLCHECK(ncclSend(((char*)sendbuff)+r*rankOffset, recvcount, datatype, r, comm, stream));
|
|
}
|
|
NCCLCHECK(ncclRecv(recvbuff, recvcount, datatype, root, comm, stream));
|
|
NCCLCHECK(ncclGroupEnd());
|
|
return ncclSuccess;
|
|
}
|
|
|
|
struct NvtxParamsSendRecv {
|
|
size_t bytes;
|
|
int peer;
|
|
};
|
|
constexpr const nvtxPayloadSchemaEntry_t SendRecvSchema[] = {
|
|
{0, NVTX_PAYLOAD_ENTRY_TYPE_SIZE, "Bytes"},
|
|
{0, NVTX_PAYLOAD_ENTRY_TYPE_INT, "Peer rank", nullptr, 0, offsetof(NvtxParamsSendRecv, peer)}
|
|
};
|
|
|
|
NCCL_API(ncclResult_t, ncclSend, const void* sendbuff, size_t count, ncclDataType_t datatype, int peer,
|
|
ncclComm_t comm, cudaStream_t stream);
|
|
ncclResult_t ncclSend(const void* sendbuff, size_t count, ncclDataType_t datatype, int peer,
|
|
ncclComm_t comm, cudaStream_t stream) {
|
|
NvtxParamsSendRecv payload{count * ncclTypeSize(datatype), peer};
|
|
NVTX3_FUNC_WITH_PARAMS(Send, SendRecvSchema, payload)
|
|
|
|
if (mscclAvailable() && !mscclIsCaller()) {
|
|
return mscclEnqueueCheck(
|
|
sendbuff, nullptr, nullptr, nullptr, nullptr, nullptr,
|
|
count, datatype, 0, peer, ncclSum, mscclFuncSend, comm, stream);
|
|
}
|
|
|
|
struct ncclInfo info = { ncclFuncSend, "Send",
|
|
NULL, (void*)sendbuff, count, datatype, ncclSum, peer, comm, stream, /* Args */
|
|
1, 1 };
|
|
ncclResult_t ret;
|
|
NCCLCHECK(ncclGroupStart());
|
|
ret = ncclEnqueueCheck(&info);
|
|
NCCLCHECK(ncclGroupEnd());
|
|
return ret;
|
|
}
|
|
|
|
NCCL_API(ncclResult_t, ncclRecv, void* recvbuff, size_t count, ncclDataType_t datatype, int peer,
|
|
ncclComm_t comm, cudaStream_t stream);
|
|
ncclResult_t ncclRecv(void* recvbuff, size_t count, ncclDataType_t datatype, int peer,
|
|
ncclComm_t comm, cudaStream_t stream) {
|
|
NvtxParamsSendRecv payload{count * ncclTypeSize(datatype), peer};
|
|
NVTX3_FUNC_WITH_PARAMS(Recv, SendRecvSchema, payload)
|
|
|
|
if (mscclAvailable() && !mscclIsCaller()) {
|
|
return mscclEnqueueCheck(
|
|
nullptr, nullptr, nullptr, recvbuff, nullptr, nullptr,
|
|
count, datatype, 0, peer, ncclSum, mscclFuncRecv, comm, stream);
|
|
}
|
|
|
|
struct ncclInfo info = { ncclFuncRecv, "Recv",
|
|
NULL, recvbuff, count, datatype, ncclSum, peer, comm, stream, /* Args */
|
|
1, 1 };
|
|
ncclResult_t ret;
|
|
NCCLCHECK(ncclGroupStart());
|
|
ret = ncclEnqueueCheck(&info);
|
|
NCCLCHECK(ncclGroupEnd());
|
|
return ret;
|
|
}
|