Merge remote-tracking branch 'nccl/master' into develop
[ROCm/rccl commit: 4f0e223db4]
Este cometimento está contido em:
Wenkai Du
@@ -12,16 +12,17 @@
|
||||
#include <hip/hip_fp16.h>
|
||||
|
||||
#define NCCL_MAJOR 2
|
||||
#define NCCL_MINOR 11
|
||||
#define NCCL_PATCH 4
|
||||
#define NCCL_MINOR 14
|
||||
#define NCCL_PATCH 3
|
||||
#define NCCL_SUFFIX ""
|
||||
|
||||
#define NCCL_VERSION_CODE 21104
|
||||
#define NCCL_VERSION_CODE 21403
|
||||
#define NCCL_VERSION(X,Y,Z) (((X) <= 2 && (Y) <= 8) ? (X) * 1000 + (Y) * 100 + (Z) : (X) * 10000 + (Y) * 100 + (Z))
|
||||
|
||||
#define RCCL_BFLOAT16 1
|
||||
#define RCCL_GATHER_SCATTER 1
|
||||
#define RCCL_ALLTOALLV 1
|
||||
#define RCCL_MULTIRANKPERGPU 1
|
||||
|
||||
#ifdef __cplusplus
|
||||
extern "C" {
|
||||
@@ -40,7 +41,29 @@ typedef enum { ncclSuccess = 0,
|
||||
ncclInternalError = 3,
|
||||
ncclInvalidArgument = 4,
|
||||
ncclInvalidUsage = 5,
|
||||
ncclNumResults = 6 } ncclResult_t;
|
||||
ncclRemoteError = 6,
|
||||
ncclInProgress = 7,
|
||||
ncclNumResults = 8 } ncclResult_t;
|
||||
|
||||
/* Communicator configuration. Users can assign value to attributes to specify the
|
||||
* behavior of a communicator. */
|
||||
typedef struct ncclConfig_v21400 {
|
||||
/* attributes that users should never touch. */
|
||||
size_t size;
|
||||
unsigned int magic;
|
||||
unsigned int version;
|
||||
/* attributes that users are able to customize. */
|
||||
int blocking;
|
||||
} ncclConfig_t;
|
||||
|
||||
/* Config initializer must be assigned to initialize config structure when it is created.
|
||||
* Not initialized config will result in NCCL error. */
|
||||
#define NCCL_CONFIG_INITIALIZER { \
|
||||
sizeof(ncclConfig_t), /* size */ \
|
||||
0xcafebeef, /* magic */ \
|
||||
NCCL_VERSION(NCCL_MAJOR, NCCL_MINOR, NCCL_PATCH), /* version */ \
|
||||
1 /* blocking */ \
|
||||
}
|
||||
|
||||
/*! @brief Return the NCCL_VERSION_CODE of the NCCL library in the supplied integer.
|
||||
*
|
||||
@@ -69,6 +92,13 @@ ncclResult_t ncclGetUniqueId(ncclUniqueId* uniqueId);
|
||||
ncclResult_t pncclGetUniqueId(ncclUniqueId* uniqueId);
|
||||
/// @endcond
|
||||
|
||||
/*! @brief Create a new communicator (multi thread/process version) with a configuration
|
||||
* set by users. */
|
||||
ncclResult_t ncclCommInitRankConfig(ncclComm_t* comm, int nranks, ncclUniqueId commId, int rank, ncclConfig_t* config);
|
||||
/// @cond include_hidden
|
||||
ncclResult_t pncclCommInitRankConfig(ncclComm_t* comm, int nranks, ncclUniqueId commId, int rank, ncclConfig_t* config);
|
||||
/// @endcond
|
||||
|
||||
/*! @brief Creates a new communicator (multi thread/process version).
|
||||
|
||||
@details
|
||||
@@ -87,6 +117,28 @@ ncclResult_t ncclCommInitRank(ncclComm_t* comm, int nranks, ncclUniqueId commId
|
||||
ncclResult_t pncclCommInitRank(ncclComm_t* comm, int nranks, ncclUniqueId commId, int rank);
|
||||
/// @endcond
|
||||
|
||||
/*! @brief Creates a new communicator (multi thread/process version) allowing multiple ranks per device.
|
||||
|
||||
@details
|
||||
rank must be between 0 and nranks-1 and unique within a communicator clique.
|
||||
Each rank is associated to a HIP device, which has to be set before calling
|
||||
ncclCommInitRankMulti.
|
||||
Since this version of the function allows multiple ranks to utilize the same
|
||||
HIP device, a unique virtualId per device has to be provided by each calling
|
||||
rank.
|
||||
ncclCommInitRankMulti implicitly syncronizes with other ranks, so it must be
|
||||
called by different threads/processes or use ncclGroupStart/ncclGroupEnd.
|
||||
|
||||
@param[in]
|
||||
comm ncclComm_t*
|
||||
communicator struct pointer
|
||||
*/
|
||||
ncclResult_t ncclCommInitRankMulti(ncclComm_t* comm, int nranks, ncclUniqueId commId, int rank, int virtualId);
|
||||
/// @cond include_hidden
|
||||
ncclResult_t pncclCommInitRankMulti(ncclComm_t* comm, int nranks, ncclUniqueId commId, int rank, int virtualId);
|
||||
/// @endcond
|
||||
|
||||
|
||||
/*! @brief Creates a clique of communicators (single process version).
|
||||
*
|
||||
* @details This is a convenience function to create a single-process communicator clique.
|
||||
@@ -100,23 +152,46 @@ ncclResult_t ncclCommInitAll(ncclComm_t* comm, int ndev, const int* devlist);
|
||||
ncclResult_t pncclCommInitAll(ncclComm_t* comm, int ndev, const int* devlist);
|
||||
/// @endcond
|
||||
|
||||
/*! @brief Frees resources associated with communicator object, but waits for any operations that might still be running on the device */
|
||||
/*! @brief Finalize a communicator.
|
||||
* @details ncclCommFinalize flushes all issued communications,
|
||||
* and marks communicator state as ncclInProgress. The state will change to ncclSuccess
|
||||
* when the communicator is globally quiescent and related resources are freed; then,
|
||||
* calling ncclCommDestroy can locally free the rest of the resources (e.g. communicator
|
||||
* itself) without blocking. */
|
||||
ncclResult_t ncclCommFinalize(ncclComm_t comm);
|
||||
/// @cond include_hidden
|
||||
ncclResult_t pncclCommFinalize(ncclComm_t comm);
|
||||
/// @endcond
|
||||
|
||||
/*! @brief Frees local resources associated with communicator object. */
|
||||
|
||||
ncclResult_t ncclCommDestroy(ncclComm_t comm);
|
||||
/// @cond include_hidden
|
||||
ncclResult_t pncclCommDestroy(ncclComm_t comm);
|
||||
/// @endcond
|
||||
|
||||
/*! @brief Frees resources associated with communicator object and aborts any operations that might still be running on the device. */
|
||||
/*! @brief Frees resources associated with communicator object and aborts any operations
|
||||
* that might still be running on the device. */
|
||||
ncclResult_t ncclCommAbort(ncclComm_t comm);
|
||||
/// @cond include_hidden
|
||||
ncclResult_t pncclCommAbort(ncclComm_t comm);
|
||||
/// @endcond
|
||||
|
||||
/*! @brief Returns a human-readable error message. */
|
||||
/*! @brief Returns a string for each error code. */
|
||||
const char* ncclGetErrorString(ncclResult_t result);
|
||||
/// @cond include_hidden
|
||||
const char* pncclGetErrorString(ncclResult_t result);
|
||||
/// @endcond
|
||||
|
||||
/*! @brief Checks whether the comm has encountered any asynchronous errors */
|
||||
/*! @brief Returns a human-readable message of the last error that occurred.
|
||||
* comm is currently unused and can be set to NULL
|
||||
*/
|
||||
const char* ncclGetLastError(ncclComm_t comm);
|
||||
/// @cond include_hidden
|
||||
const char* pncclGetError(ncclComm_t comm);
|
||||
/// @endcond
|
||||
|
||||
/* Checks whether the comm has encountered any asynchronous errors */
|
||||
ncclResult_t ncclCommGetAsyncError(ncclComm_t comm, ncclResult_t *asyncError);
|
||||
/// @cond include_hidden
|
||||
ncclResult_t pncclCommGetAsyncError(ncclComm_t comm, ncclResult_t *asyncError);
|
||||
@@ -173,7 +248,7 @@ typedef enum { ncclInt8 = 0, ncclChar = 0,
|
||||
ncclBfloat16 = 9,
|
||||
ncclNumTypes = 10 } ncclDataType_t;
|
||||
|
||||
/* ncclScalarResidence_t: Location and dereferencing logic for scalar arguments. */
|
||||
/*! @brief ncclScalarResidence_t: Location and dereferencing logic for scalar arguments. */
|
||||
typedef enum {
|
||||
/* ncclScalarDevice: The scalar is in device-visible memory and will be
|
||||
* dereferenced while the collective is running. */
|
||||
@@ -184,9 +259,7 @@ typedef enum {
|
||||
ncclScalarHostImmediate = 1
|
||||
} ncclScalarResidence_t;
|
||||
|
||||
/*
|
||||
* ncclRedOpCreatePreMulSum
|
||||
*
|
||||
/*! @brief ncclRedOpCreatePreMulSum
|
||||
* Creates a new reduction operator which pre-multiplies input values by a given
|
||||
* scalar locally before reducing them with peer values via summation. For use
|
||||
* only with collectives launched against *comm* and *datatype*. The
|
||||
@@ -195,17 +268,19 @@ typedef enum {
|
||||
* is stored in *op*.
|
||||
*/
|
||||
ncclResult_t ncclRedOpCreatePreMulSum(ncclRedOp_t *op, void *scalar, ncclDataType_t datatype, ncclScalarResidence_t residence, ncclComm_t comm);
|
||||
/// @cond include_hidden
|
||||
ncclResult_t pncclRedOpCreatePreMulSum(ncclRedOp_t *op, void *scalar, ncclDataType_t datatype, ncclScalarResidence_t residence, ncclComm_t comm);
|
||||
/// @endcond
|
||||
|
||||
/*
|
||||
* ncclRedOpDestroy
|
||||
*
|
||||
* Destroys the reduction operator *op*. The operator must have been created by
|
||||
/*! @brief ncclRedOpDestroy
|
||||
* @details Destroys the reduction operator *op*. The operator must have been created by
|
||||
* ncclRedOpCreatePreMul with the matching communicator *comm*. An operator may be
|
||||
* destroyed as soon as the last NCCL function which is given that operator returns.
|
||||
*/
|
||||
ncclResult_t ncclRedOpDestroy(ncclRedOp_t op, ncclComm_t comm);
|
||||
/// @cond include_hidden
|
||||
ncclResult_t pncclRedOpDestroy(ncclRedOp_t op, ncclComm_t comm);
|
||||
/// @endcond
|
||||
|
||||
/*
|
||||
* Collective communication operations
|
||||
@@ -345,11 +420,11 @@ ncclResult_t pncclSend(const void* sendbuff, size_t count, ncclDataType_t dataty
|
||||
* need to progress concurrently to complete, they must be fused within a ncclGroupStart/
|
||||
* ncclGroupEnd section.
|
||||
*/
|
||||
ncclResult_t ncclRecv(void* recvbuff, size_t count, ncclDataType_t datatype, int peer,
|
||||
ncclComm_t comm, hipStream_t stream);
|
||||
/// @cond include_hidden
|
||||
ncclResult_t pncclRecv(void* recvbuff, size_t count, ncclDataType_t datatype, int peer,
|
||||
ncclComm_t comm, hipStream_t stream);
|
||||
ncclResult_t ncclRecv(void* recvbuff, size_t count, ncclDataType_t datatype, int peer,
|
||||
ncclComm_t comm, hipStream_t stream);
|
||||
/// @endcond
|
||||
|
||||
/*! @brief Gather
|
||||
|
||||
@@ -13,8 +13,8 @@ struct ncclGraphInfo {
|
||||
int pattern;
|
||||
int nChannels;
|
||||
int sameChannels;
|
||||
float speedIntra;
|
||||
float speedInter;
|
||||
float bwIntra;
|
||||
float bwInter;
|
||||
int typeIntra;
|
||||
int typeInter;
|
||||
};
|
||||
@@ -28,6 +28,7 @@ struct allGather3Data_t{
|
||||
struct ncclGraphInfo collNet;
|
||||
struct ncclTopoRanks topoRanks;
|
||||
bool pivotA2AEnabled;
|
||||
bool ll128Enabled;
|
||||
};
|
||||
|
||||
void initCollNet();
|
||||
|
||||
@@ -75,14 +75,39 @@ int busIdToCudaDev(int64_t busId) {
|
||||
return node_model->busIdToCudaDev(busId);
|
||||
}
|
||||
|
||||
static int useMemcpy = 0;
|
||||
|
||||
/* Determine if two peers can communicate with P2P */
|
||||
ncclResult_t p2pCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTopoGraph* graph, struct ncclPeerInfo* info1, struct ncclPeerInfo* info2) {
|
||||
// Rule out different nodes
|
||||
*ret = 0;
|
||||
if (info1->hostHash != info2->hostHash) return ncclSuccess;
|
||||
int cudaDev1 = busIdToCudaDev(info1->busId);
|
||||
int cudaDev2 = busIdToCudaDev(info2->busId);
|
||||
*ret = node_model->p2pCanConnect(cudaDev1, cudaDev2);
|
||||
if (!info1->hasFineGrain || !info2->hasFineGrain) {
|
||||
*ret = 0;
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
// Rule out different nodes / isolated containers
|
||||
if (info1->hostHash != info2->hostHash || info1->shmDev != info2->shmDev) {
|
||||
*ret = 0;
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
// Check topology / p2p level.
|
||||
int intermediateRank;
|
||||
NCCLCHECK(ncclTopoCheckP2p(topo, info1->busId, info2->busId, ret, NULL, &intermediateRank));
|
||||
if (*ret == 0) return ncclSuccess;
|
||||
if (intermediateRank != -1) {
|
||||
if (useMemcpy) *ret = 0;
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
// Check if NET would work better
|
||||
int useNet = 0;
|
||||
NCCLCHECK(ncclTopoCheckNet(topo, info1->busId, info2->busId, &useNet));
|
||||
if (useNet) {
|
||||
*ret = 0;
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
*ret = 1;
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
@@ -118,14 +143,26 @@ struct ncclTransport p2pTransport = {
|
||||
{ p2pRecvSetup, NULL, NULL, NULL }
|
||||
};
|
||||
|
||||
NCCL_PARAM(ShmDisable, "SHM_DISABLE", 0);
|
||||
|
||||
/* Determine if two peers can communicate with SHM */
|
||||
ncclResult_t shmCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTopoGraph* graph, struct ncclPeerInfo* info1, struct ncclPeerInfo* info2) {
|
||||
// Rule out different nodes
|
||||
*ret = 0;
|
||||
if (ncclParamShmDisable() == 1) return ncclSuccess;
|
||||
|
||||
int useNet = 0;
|
||||
NCCLCHECK(ncclTopoCheckNet(topo, info1->busId, info2->busId, &useNet));
|
||||
if (useNet) return ncclSuccess;
|
||||
|
||||
// Same host?
|
||||
TRACE(NCCL_INIT|NCCL_SHM, "peer1 hostHash %lx peer2 hostHash %lx", info1->hostHash, info2->hostHash);
|
||||
if (info1->hostHash != info2->hostHash) return ncclSuccess;
|
||||
int cudaDev1 = busIdToCudaDev(info1->busId);
|
||||
int cudaDev2 = busIdToCudaDev(info2->busId);
|
||||
*ret = node_model->shmCanConnect(cudaDev1, cudaDev2);
|
||||
|
||||
// Common /dev/shm (between containers) ?
|
||||
TRACE(NCCL_INIT|NCCL_SHM, "peer1 shmDev %lx peer2 shmDev %lx", info1->shmDev, info2->shmDev);
|
||||
if (info1->shmDev != info2->shmDev) return ncclSuccess;
|
||||
|
||||
*ret = 1;
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
@@ -161,7 +198,7 @@ struct setupReq {
|
||||
|
||||
/* Determine if two peers can communicate with NET */
|
||||
ncclResult_t netCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTopoGraph* graph, struct ncclPeerInfo* info1, struct ncclPeerInfo* info2) {
|
||||
*ret = node_model->netCanConnect(info1->rank, info2->rank);
|
||||
*ret = 1;
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
|
||||
@@ -224,6 +224,7 @@ int main(int argc,char* argv[])
|
||||
comm[i].topo = node_model->getSystem(i);
|
||||
comm[i].peerInfo = peerInfo;
|
||||
comm[i].ncclNet = ncclNet;
|
||||
comm[i].virtualId = -1;
|
||||
// Mark channels as non initialized.
|
||||
for (int c=0; c<MAXCHANNELS; c++) comm[i].channels[c].id = -1;
|
||||
NCCLCHECK(fillInfo(&comm[i], comm[i].peerInfo+comm[i].rank, 0));
|
||||
|
||||
@@ -39,7 +39,6 @@ extern NodeModel *node_model;
|
||||
|
||||
NCCL_PARAM(CollNetEnable, "COLLNET_ENABLE", 0);
|
||||
NCCL_PARAM(GraphDumpFileRank, "GRAPH_DUMP_FILE_RANK", 0);
|
||||
NCCL_PARAM(CollNetNodeThreshold, "COLLNET_NODE_THRESHOLD", 2);
|
||||
|
||||
thread_local int ncclDebugNoWarn = 0;
|
||||
ncclCollNet_t* ncclCollNet = NULL;
|
||||
@@ -462,7 +461,10 @@ ncclResult_t ncclTransportCollNetFree(struct ncclComm* comm) {
|
||||
}
|
||||
|
||||
RCCL_PARAM(P2pNetDisable, "P2P_NET_DISABLE", 0);
|
||||
NCCL_PARAM(CollNetNodeThreshold, "COLLNET_NODE_THRESHOLD", 2);
|
||||
RCCL_PARAM(PivotAlltoallEnable, "PIVOT_ALLTOALL_ENABLE", 0);
|
||||
NCCL_PARAM(AllocP2pNetLLBuffers, "NCCL_ALLOC_P2P_NET_LL_BUFFERS", 0);
|
||||
RCCL_PARAM(LL128ForceEnable, "LL128_FORCE_ENABLE", 0);
|
||||
|
||||
ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather3Data_t *allGather3Data,
|
||||
struct ncclTopoGraph& treeGraph, struct ncclTopoGraph& ringGraph, struct ncclTopoGraph& collNetGraph) {
|
||||
@@ -484,13 +486,30 @@ ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather3Data_t
|
||||
//NCCLCHECK(fillInfo(comm, comm->peerInfo+rank, comm->rank));
|
||||
//NCCLCHECK(bootstrapAllGather(comm->bootstrap, comm->peerInfo, sizeof(struct ncclPeerInfo)));
|
||||
|
||||
for (int i = 0; i < nranks; i++) {
|
||||
if ((i != rank) && (comm->peerInfo[i].hostHash == comm->peerInfo[rank].hostHash) && (comm->peerInfo[i].busId == comm->peerInfo[rank].busId)) {
|
||||
WARN("Duplicate GPU detected : rank %d and rank %d both on CUDA device %lx", rank, i, comm->peerInfo[rank].busId);
|
||||
return ncclInvalidUsage;
|
||||
//If virtualId == -1 multiRank support has not been requested by user, using original interface
|
||||
if (comm->virtualId == -1) {
|
||||
for (int i = 0; i < nranks; i++) {
|
||||
if ((i != rank) && (comm->peerInfo[i].hostHash == comm->peerInfo[rank].hostHash) && (comm->peerInfo[i].busId == comm->peerInfo[rank].busId)) {
|
||||
WARN("Duplicate GPU detected : rank %d and rank %d both on CUDA device %lx", rank, i, comm->peerInfo[rank].busId);
|
||||
return ncclInvalidUsage;
|
||||
}
|
||||
}
|
||||
}
|
||||
else {
|
||||
//Multiple ranks can use the same device, but need to have different virtualId's.
|
||||
for (int i = 0; i < nranks; i++) {
|
||||
for (int j=0; j < nranks; j++) {
|
||||
if (j==i) continue;
|
||||
if((comm->peerInfo[i].hostHash == comm->peerInfo[j].hostHash) &&
|
||||
(comm->peerInfo[i].busId == comm->peerInfo[j].busId) &&
|
||||
(comm->peerInfo[i].virtualId == comm->peerInfo[j].virtualId)) {
|
||||
WARN("Duplicate virtualId detected : rank %d and rank %d both on GPU device %lx virtualId %d",
|
||||
i, j, comm->peerInfo[rank].busId, comm->peerInfo[i].virtualId);
|
||||
return ncclInvalidUsage;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// AllGather1 - end
|
||||
|
||||
// Topo detection / System graph creation
|
||||
@@ -502,6 +521,8 @@ ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather3Data_t
|
||||
// init Pivot A2A related fields
|
||||
comm->topo->pivotA2AEnabled = false;
|
||||
comm->topo->pivotA2ANumBiRings = 0;
|
||||
// LL128
|
||||
comm->topo->ll128Enabled = false;
|
||||
// Compute paths between GPUs and NICs
|
||||
NCCLCHECK(ncclTopoComputePaths(comm->topo, comm));
|
||||
// Remove inaccessible GPUs and unused NICs
|
||||
@@ -573,6 +594,8 @@ ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather3Data_t
|
||||
allXgmi &= isXGMI;
|
||||
}
|
||||
}
|
||||
// Initialize num P2P LL buffers for this communicator
|
||||
comm->allocP2pNetLLBuffers = ncclParamAllocP2pNetLLBuffers() == 1;
|
||||
|
||||
if (comm->rank == ncclParamGraphDumpFileRank()) {
|
||||
struct ncclTopoGraph* graphs[3] = { &ringGraph, &treeGraph, &collNetGraph };
|
||||
@@ -605,8 +628,8 @@ ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather3Data_t
|
||||
int pattern;
|
||||
int nChannels;
|
||||
int sameChannels;
|
||||
float speedIntra;
|
||||
float speedInter;
|
||||
float bwIntra;
|
||||
float bwInter;
|
||||
int typeIntra;
|
||||
int typeInter;
|
||||
};
|
||||
@@ -620,6 +643,7 @@ ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather3Data_t
|
||||
struct ncclGraphInfo collNet;
|
||||
struct ncclTopoRanks topoRanks;
|
||||
bool pivotA2AEnabled;
|
||||
bool ll128Enabled;
|
||||
} *allGather3Data;
|
||||
|
||||
NCCLCHECK(ncclCalloc(&allGather3Data, nranks));
|
||||
@@ -627,13 +651,19 @@ ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather3Data_t
|
||||
int idx;
|
||||
NCCLCHECK(ncclTopoIdToIndex(comm->topo, GPU, comm->busId, &idx));
|
||||
allGather3Data[rank].nc = 2;
|
||||
if (comm->topo->nodes[GPU].count == comm->topo->nRanks && comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 906 && allXgmi)
|
||||
if ( ((comm->topo->nodes[GPU].count == comm->topo->nRanks && comm->virtualId == -1) ||
|
||||
(comm->topo->nodes[GPU].count <= comm->topo->nRanks && comm->virtualId != -1)) &&
|
||||
comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 906 && allXgmi)
|
||||
allGather3Data[rank].nc = 4;
|
||||
if (comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 908)
|
||||
allGather3Data[rank].nc = std::max(4/ringGraph.nChannels, 2);
|
||||
if (comm->topo->nodes[GPU].count == comm->topo->nRanks && (comm->topo->type & RCCL_TOPO_CR8G))
|
||||
if ( ((comm->topo->nodes[GPU].count == comm->topo->nRanks && comm->virtualId == -1) ||
|
||||
(comm->topo->nodes[GPU].count <= comm->topo->nRanks && comm->virtualId != -1)) &&
|
||||
(comm->topo->type & RCCL_TOPO_CR8G))
|
||||
allGather3Data[rank].nc = 4;
|
||||
if (comm->topo->nodes[GPU].count == comm->topo->nRanks && comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 910)
|
||||
if (((comm->topo->nodes[GPU].count == comm->topo->nRanks && comm->virtualId == -1) ||
|
||||
(comm->topo->nodes[GPU].count <= comm->topo->nRanks && comm->virtualId != -1)) &&
|
||||
comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 910)
|
||||
allGather3Data[rank].nc = 4;
|
||||
if (comm->topo->nodes[GPU].nodes[idx].gpu.gcn == 910)
|
||||
allGather3Data[rank].nc = std::max(allGather3Data[rank].nc, 4/ringGraph.nChannels);
|
||||
@@ -643,26 +673,28 @@ ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather3Data_t
|
||||
allGather3Data[rank].tree.pattern = treeGraph.pattern;
|
||||
allGather3Data[rank].tree.nChannels = treeGraph.nChannels;
|
||||
allGather3Data[rank].tree.sameChannels = treeGraph.sameChannels;
|
||||
allGather3Data[rank].tree.speedIntra = treeGraph.speedIntra;
|
||||
allGather3Data[rank].tree.speedInter = treeGraph.speedInter;
|
||||
allGather3Data[rank].tree.bwIntra = treeGraph.bwIntra;
|
||||
allGather3Data[rank].tree.bwInter = treeGraph.bwInter;
|
||||
allGather3Data[rank].tree.typeIntra = treeGraph.typeIntra;
|
||||
allGather3Data[rank].tree.typeInter = treeGraph.typeInter;
|
||||
allGather3Data[rank].ring.pattern = ringGraph.pattern;
|
||||
allGather3Data[rank].ring.nChannels = ringGraph.nChannels;
|
||||
allGather3Data[rank].ring.sameChannels = ringGraph.sameChannels;
|
||||
allGather3Data[rank].ring.speedIntra = ringGraph.speedIntra;
|
||||
allGather3Data[rank].ring.speedInter = ringGraph.speedInter;
|
||||
allGather3Data[rank].ring.bwIntra = ringGraph.bwIntra;
|
||||
allGather3Data[rank].ring.bwInter = ringGraph.bwInter;
|
||||
allGather3Data[rank].ring.typeIntra = ringGraph.typeIntra;
|
||||
allGather3Data[rank].ring.typeInter = ringGraph.typeInter;
|
||||
allGather3Data[rank].collNet.pattern = collNetGraph.pattern;
|
||||
allGather3Data[rank].collNet.nChannels = collNetGraph.nChannels;
|
||||
allGather3Data[rank].collNet.sameChannels = collNetGraph.sameChannels;
|
||||
allGather3Data[rank].collNet.speedIntra = collNetGraph.speedIntra;
|
||||
allGather3Data[rank].collNet.speedInter = collNetGraph.speedInter;
|
||||
allGather3Data[rank].collNet.bwIntra = collNetGraph.bwIntra;
|
||||
allGather3Data[rank].collNet.bwInter = collNetGraph.bwInter;
|
||||
allGather3Data[rank].collNet.typeIntra = collNetGraph.typeIntra;
|
||||
allGather3Data[rank].collNet.typeInter = collNetGraph.typeInter;
|
||||
allGather3Data[rank].collNetSupport = comm->collNetSupport;
|
||||
allGather3Data[rank].pivotA2AEnabled = comm->topo->pivotA2AEnabled && rcclParamPivotAlltoallEnable();
|
||||
comm->topo->ll128Enabled = comm->topo->ll128Enabled || rcclParamLL128ForceEnable();
|
||||
allGather3Data[rank].ll128Enabled = comm->topo->ll128Enabled;
|
||||
|
||||
comm->nChannels = (comm->topo->nodes[GPU].count != comm->topo->nRanks && comm->topo->nodes[NET].count)
|
||||
? std::min(treeGraph.nChannels, ringGraph.nChannels) : ringGraph.nChannels;
|
||||
@@ -738,24 +770,25 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t
|
||||
// Make sure we align all ranks so that the tuning is consistent across ranks
|
||||
treeGraph.nChannels = std::min(allGather3Data[i].tree.nChannels, treeGraph.nChannels);
|
||||
treeGraph.sameChannels = std::min(allGather3Data[i].tree.sameChannels, treeGraph.sameChannels);
|
||||
treeGraph.speedIntra = std::min(allGather3Data[i].tree.speedIntra, treeGraph.speedIntra);
|
||||
treeGraph.speedInter = std::min(allGather3Data[i].tree.speedInter, treeGraph.speedInter);
|
||||
treeGraph.bwIntra = std::min(allGather3Data[i].tree.bwIntra, treeGraph.bwIntra);
|
||||
treeGraph.bwInter = std::min(allGather3Data[i].tree.bwInter, treeGraph.bwInter);
|
||||
treeGraph.typeIntra = std::max(allGather3Data[i].tree.typeIntra, treeGraph.typeIntra);
|
||||
treeGraph.typeInter = std::max(allGather3Data[i].tree.typeInter, treeGraph.typeInter);
|
||||
ringGraph.nChannels = std::min(allGather3Data[i].ring.nChannels, ringGraph.nChannels);
|
||||
ringGraph.sameChannels = std::min(allGather3Data[i].ring.sameChannels, ringGraph.sameChannels);
|
||||
ringGraph.speedIntra = std::min(allGather3Data[i].ring.speedIntra, ringGraph.speedIntra);
|
||||
ringGraph.speedInter = std::min(allGather3Data[i].ring.speedInter, ringGraph.speedInter);
|
||||
ringGraph.bwIntra = std::min(allGather3Data[i].ring.bwIntra, ringGraph.bwIntra);
|
||||
ringGraph.bwInter = std::min(allGather3Data[i].ring.bwInter, ringGraph.bwInter);
|
||||
ringGraph.typeIntra = std::max(allGather3Data[i].ring.typeIntra, ringGraph.typeIntra);
|
||||
ringGraph.typeInter = std::max(allGather3Data[i].ring.typeInter, ringGraph.typeInter);
|
||||
collNetGraph.nChannels = std::min(allGather3Data[i].collNet.nChannels, collNetGraph.nChannels);
|
||||
collNetGraph.sameChannels = std::min(allGather3Data[i].collNet.sameChannels, collNetGraph.sameChannels);
|
||||
collNetGraph.speedIntra = std::min(allGather3Data[i].collNet.speedIntra, collNetGraph.speedIntra);
|
||||
collNetGraph.speedInter = std::min(allGather3Data[i].collNet.speedInter, collNetGraph.speedInter);
|
||||
collNetGraph.bwIntra = std::min(allGather3Data[i].collNet.bwIntra, collNetGraph.bwIntra);
|
||||
collNetGraph.bwInter = std::min(allGather3Data[i].collNet.bwInter, collNetGraph.bwInter);
|
||||
collNetGraph.typeIntra = std::max(allGather3Data[i].collNet.typeIntra, collNetGraph.typeIntra);
|
||||
collNetGraph.typeInter = std::max(allGather3Data[i].collNet.typeInter, collNetGraph.typeInter);
|
||||
comm->collNetSupport = std::min(allGather3Data[i].collNetSupport, comm->collNetSupport);
|
||||
comm->topo->pivotA2AEnabled = comm->topo->pivotA2AEnabled && allGather3Data[i].pivotA2AEnabled;
|
||||
comm->topo->ll128Enabled = comm->topo->ll128Enabled && allGather3Data[i].ll128Enabled;
|
||||
}
|
||||
|
||||
comm->nChannels = treeGraph.nChannels = ringGraph.nChannels =
|
||||
@@ -787,7 +820,10 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t
|
||||
NCCLCHECK(ncclCalloc(&rings, nranks*MAXCHANNELS));
|
||||
NCCLCHECK(ncclTopoPostset(comm, nodesFirstRank, nodesTreePatterns, allTopoRanks, rings, &collNetGraph, nc));
|
||||
|
||||
if (comm->topo->pivotA2ANumBiRings == 3) NCCLCHECK(ncclTreeBasePostset(comm, &treeGraph));
|
||||
if (comm->topo->pivotA2ANumBiRings == 3) {
|
||||
NCCLCHECK(ncclTreeBasePostset(comm, &treeGraph));
|
||||
NCCLCHECK(ncclBinaryTreePostset(comm, &treeGraph));
|
||||
}
|
||||
|
||||
free(allTopoRanks);
|
||||
free(nodesTreePatterns);
|
||||
@@ -798,16 +834,25 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t
|
||||
|
||||
TRACE(NCCL_INIT, "rank %d nranks %d - BUILT %d TREES/RINGS", rank, nranks, comm->nChannels);
|
||||
|
||||
char line[1024];
|
||||
char line[1024], binline[1024];
|
||||
line[0]='\0';
|
||||
binline[0]='\0';
|
||||
for (int c=0; c<comm->nChannels; c++) {
|
||||
struct ncclTree* tree = &comm->channels[c].tree;
|
||||
struct ncclTree* binTree = &comm->channels[c].binTree;
|
||||
snprintf(line+strlen(line), 1023-strlen(line), " [%d] %d/%d/%d->%d->%d",
|
||||
c, tree->down[0], tree->down[1], tree->down[2], rank, tree->up);
|
||||
if (comm->topo->pivotA2ANumBiRings == 3)
|
||||
snprintf(binline+strlen(binline), 1023-strlen(binline), " [%d] %d/%d/%d->%d->%d",
|
||||
c, binTree->down[0], binTree->down[1], binTree->down[2], rank, binTree->up);
|
||||
INFO(NCCL_GRAPH, "Ring %d : %d -> %d -> %d", c, comm->channels[c].ring.prev, comm->rank, comm->channels[c].ring.next);
|
||||
}
|
||||
line[1023] = '\0';
|
||||
INFO(NCCL_INIT, "Trees%s", line);
|
||||
if (comm->topo->pivotA2ANumBiRings == 3) {
|
||||
binline[1023] = '\0';
|
||||
INFO(NCCL_INIT, "BinTrees%s", binline);
|
||||
}
|
||||
|
||||
//NCCLCHECK(computeBuffSizes(comm));
|
||||
|
||||
@@ -838,6 +883,11 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t
|
||||
if (comm->nRanks == 1) continue;
|
||||
NCCLCHECKGOTO(ncclTransportP2pConnect(comm, c, NCCL_MAX_TREE_ARITY, channel->tree.down, 1, &channel->tree.up, 0), ret, affinity_restore);
|
||||
NCCLCHECKGOTO(ncclTransportP2pConnect(comm, c, 1, &channel->tree.up, NCCL_MAX_TREE_ARITY, channel->tree.down, 0), ret, affinity_restore);
|
||||
// RCCL: need to connect binTree as well
|
||||
if (comm->topo->pivotA2ANumBiRings == 3) {
|
||||
NCCLCHECKGOTO(ncclTransportP2pConnect(comm, c, NCCL_MAX_TREE_ARITY, channel->binTree.down, 1, &channel->binTree.up, 0), ret, affinity_restore);
|
||||
NCCLCHECKGOTO(ncclTransportP2pConnect(comm, c, 1, &channel->binTree.up, NCCL_MAX_TREE_ARITY, channel->binTree.down, 0), ret, affinity_restore);
|
||||
}
|
||||
}
|
||||
NCCLCHECKGOTO(ncclTransportP2pSetup(comm, &treeGraph, 0), ret, affinity_restore);
|
||||
INFO(NCCL_INIT, "Connected all trees");
|
||||
@@ -870,16 +920,38 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t
|
||||
NCCLCHECKGOTO(ncclTransportCollNetCheck(comm, collNetSetupFail), ret, collnet_cleanup);
|
||||
TRACE(NCCL_INIT, "rank %d Connected inter-node CollNet", rank);
|
||||
|
||||
// Connect intra-node CollNet
|
||||
char line[1024];
|
||||
line[0]='\0';
|
||||
for (int c=0; c<comm->nChannels; c++) {
|
||||
struct ncclTree* chain = &comm->channels[c].collnetChain;
|
||||
snprintf(line+strlen(line), 1023-strlen(line), " [%d] %d->%d->%d",
|
||||
c, chain->down[0], rank, chain->up);
|
||||
}
|
||||
line[1023] = '\0';
|
||||
INFO(NCCL_INIT, "Collnet Chains %s", line);
|
||||
// Connect Collnet + chain
|
||||
for (int c=0; c<comm->nChannels; c++) {
|
||||
struct ncclChannel* channel = comm->channels+c;
|
||||
NCCLCHECKGOTO(ncclTransportP2pConnect(comm, c, 1, &channel->collnetChain.up, 1, channel->collnetChain.down, 0), ret, collnet_cleanup);
|
||||
}
|
||||
NCCLCHECKGOTO(ncclTransportP2pSetup(comm, &collNetGraph, 0), ret, collnet_cleanup);
|
||||
for (int c=0; c<comm->nChannels; c++) {
|
||||
struct ncclChannel* channel = comm->channels+c;
|
||||
NCCLCHECKGOTO(ncclTransportP2pConnect(comm, c, 1, channel->collnetChain.down, 1, &channel->collnetChain.up, 1), ret, collnet_cleanup);
|
||||
}
|
||||
NCCLCHECKGOTO(ncclTransportP2pSetup(comm, &collNetGraph, 1), ret, collnet_cleanup);
|
||||
INFO(NCCL_INIT, "Connected collnet + chain");
|
||||
|
||||
// Connect intra-node CollNet + Direct
|
||||
int highestTransportType0, highestTransportType1;
|
||||
for (int c=0; c<comm->nChannels; c++) {
|
||||
struct ncclChannel* channelRecv = comm->channels+c;
|
||||
NCCLCHECKGOTO(ncclTransportP2pConnect(comm, c, NCCL_MAX_DIRECT_ARITY, channelRecv->collTree.up, NCCL_MAX_DIRECT_ARITY, channelRecv->collTree.down, 0), ret, collnet_cleanup);
|
||||
NCCLCHECKGOTO(ncclTransportP2pConnect(comm, c, NCCL_MAX_DIRECT_ARITY, channelRecv->collnetDirect.up, NCCL_MAX_DIRECT_ARITY, channelRecv->collnetDirect.down, 0), ret, collnet_cleanup);
|
||||
}
|
||||
NCCLCHECKGOTO(ncclTransportP2pSetup(comm, &collNetGraph, 0, &highestTransportType0), ret, collnet_cleanup);
|
||||
for (int c=0; c<comm->nChannels; c++) {
|
||||
struct ncclChannel* channelSend = comm->channels+c;
|
||||
NCCLCHECKGOTO(ncclTransportP2pConnect(comm, c, NCCL_MAX_DIRECT_ARITY, channelSend->collTree.down, NCCL_MAX_DIRECT_ARITY, channelSend->collTree.up, 1), ret, collnet_cleanup);
|
||||
NCCLCHECKGOTO(ncclTransportP2pConnect(comm, c, NCCL_MAX_DIRECT_ARITY, channelSend->collnetDirect.down, NCCL_MAX_DIRECT_ARITY, channelSend->collnetDirect.up, 1), ret, collnet_cleanup);
|
||||
}
|
||||
NCCLCHECKGOTO(ncclTransportP2pSetup(comm, &collNetGraph, 1, &highestTransportType1), ret, collnet_cleanup);
|
||||
|
||||
@@ -1057,6 +1129,8 @@ collnet_cleanup:
|
||||
}
|
||||
}
|
||||
|
||||
NCCLCHECKGOTO(devCommSetup(comm), ret, affinity_restore);
|
||||
|
||||
/* Local intra-node barrier */
|
||||
//NCCLCHECK(bootstrapBarrier(comm->bootstrap, comm->localRankToRank, comm->localRank, comm->localRanks, comm->localRankToRank[0]));
|
||||
|
||||
|
||||
Criar uma nova questão referindo esta
Bloquear um utilizador