Merge remote-tracking branch 'nccl/master' into develop
[ROCm/rccl commit: 3f94267f21]
This commit is contained in:
BertanDogancay
Corey Derochie
@@ -36,9 +36,8 @@ define COMPILE
|
||||
$(call COMPILE$(or $3,$(suffix $2)),$1,$2)
|
||||
endef
|
||||
|
||||
ifeq ($(shell echo "$$((1000*$(CUDA_MAJOR) + 10*$(CUDA_MINOR) >= 12080))"),1)
|
||||
NVCC_GENCODE_LDMC_FP8 = -gencode=arch=compute_100a,code=sm_100a \
|
||||
-gencode=arch=compute_120a,code=sm_120a
|
||||
ifeq ($(shell echo "$$((1000*$(CUDA_MAJOR) + 10*$(CUDA_MINOR) >= 12090))"),1)
|
||||
NVCC_GENCODE_LDMC_FP8 = -gencode=arch=compute_100f,code=sm_100f
|
||||
else ifeq ($(shell echo "$$((1000*$(CUDA_MAJOR) + 10*$(CUDA_MINOR) >= 12070))"),1)
|
||||
NVCC_GENCODE_LDMC_FP8 = -gencode=arch=compute_100a,code=sm_100a
|
||||
else
|
||||
|
||||
@@ -636,6 +636,7 @@ __device__ __forceinline__ void ncclKernelMain(struct ncclDevKernelArgs const* a
|
||||
__syncthreads();
|
||||
if (COLLTRACE && tid%WARP_SIZE == 0) traceKernelLaunch(ncclCollTraceCollLaunchType, batchIx);
|
||||
}
|
||||
profiler(FINI);
|
||||
if (COLLTRACE && tid%WARP_SIZE == 0) traceKernelEnd(ncclCollTraceKernelEndType);
|
||||
|
||||
#ifdef ENABLE_PROFILING
|
||||
|
||||
@@ -1073,7 +1073,7 @@ struct Apply_LoadMultimem {
|
||||
#endif
|
||||
|
||||
#if defined(RCCL_BFLOAT16)
|
||||
#if NCCL_CUDA_ARCH_FAMILY_SPECIFIC == 1000 || NCCL_CUDA_ARCH_FAMILY_SPECIFIC == 1010 || NCCL_CUDA_ARCH_SPECIFIC == 1200 || NCCL_CUDA_ARCH_SPECIFIC == 1210
|
||||
#if NCCL_CUDA_ARCH_SPECIFIC == 1000 || NCCL_CUDA_ARCH_SPECIFIC == 1010 || NCCL_CUDA_ARCH_FAMILY_SPECIFIC == 1000 || NCCL_CUDA_ARCH_FAMILY_SPECIFIC == 1010 || NCCL_CUDA_ARCH_SPECIFIC == 1200 || NCCL_CUDA_ARCH_SPECIFIC == 1210
|
||||
DEFINE_Apply_LoadMultimem_sum_v4_and_xparts(__nv_fp8_e4m3, e4m3x4, 4)
|
||||
DEFINE_Apply_LoadMultimem_minmax_v4_and_xparts(__nv_fp8_e4m3, e4m3x4, 4)
|
||||
DEFINE_Apply_LoadMultimem_sum_v4_and_xparts(__nv_fp8_e5m2, e5m2x4, 4)
|
||||
|
||||
@@ -108,7 +108,7 @@ def required_cuda(k):
|
||||
if k.algo in ldmc_algos:
|
||||
cudart = 12070
|
||||
arch = None
|
||||
specific_sms = [100, 120]
|
||||
specific_sms = ["100a", "101a", "100f", "101f", "120a", "121a"]
|
||||
return (cudart, arch, specific_sms)
|
||||
|
||||
################################################################################
|
||||
@@ -145,7 +145,7 @@ def kernel_conds(k):
|
||||
if not specific_sms:
|
||||
arch_cond = "__CUDA_ARCH__ >= %d"%arch
|
||||
else:
|
||||
arch_cond = " || ".join(["0"] + ["NCCL_CUDA_ARCH_SPECIFIC==%d"%(10*sm) for sm in specific_sms])
|
||||
arch_cond = " || ".join(["0"] + ["NCCL_CUDA_ARCH_%sSPECIFIC==%d"%("FAMILY_" if sm[-1] == "f" else "", 10*int(sm.replace('a', '').replace('f', ''))) for sm in specific_sms])
|
||||
return cudart_cond, arch_cond
|
||||
|
||||
def instantiate(k):
|
||||
|
||||
@@ -109,12 +109,9 @@ ncclResult_t ncclInitKernelsForDevice(int cudaArch, int maxSharedMem, size_t* ma
|
||||
if (fn == nullptr) continue;
|
||||
|
||||
cudaError_t errcode = cudaFuncGetAttributes(&attr, fn);
|
||||
if (errcode == cudaErrorNoKernelImageForDevice) continue;
|
||||
CUDACHECKGOTO(errcode, result, ignore0);
|
||||
|
||||
if (errcode != cudaSuccess) continue; // Silently ignore failures
|
||||
if (maxStackSize) {
|
||||
if (attr.localSizeBytes > *maxStackSize) *maxStackSize = attr.localSizeBytes;
|
||||
ignore0:;
|
||||
}
|
||||
if (carveout) {
|
||||
CUDACHECKGOTO(cudaFuncSetAttribute(fn,
|
||||
|
||||
@@ -180,6 +180,13 @@ ncclResult_t ncclGetLocalCpu(struct ncclTopoSystem* system, int gpu, int* retCpu
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
static int mergePathType(int type0, int type1){
|
||||
int max = std::max(type0,type1);
|
||||
int min = std::min(type0,type1);
|
||||
if(max == PATH_PHB && min == PATH_C2C) return PATH_P2C;
|
||||
else return max;
|
||||
}
|
||||
|
||||
static ncclResult_t addInterStep(struct ncclTopoSystem* system, int tx, int ix, int t1, int i1, int t2, int i2) {
|
||||
struct ncclTopoNode* cpuNode = system->nodes[tx].nodes+ix;
|
||||
struct ncclTopoNode* srcNode = system->nodes[t1].nodes+i1;
|
||||
@@ -192,7 +199,7 @@ static ncclResult_t addInterStep(struct ncclTopoSystem* system, int tx, int ix,
|
||||
|
||||
// Update path characteristics
|
||||
srcNode->paths[t2][i2].count = l;
|
||||
srcNode->paths[t2][i2].type = std::max(srcNode->paths[tx][ix].type, cpuNode->paths[t2][i2].type);
|
||||
srcNode->paths[t2][i2].type = mergePathType(srcNode->paths[tx][ix].type, cpuNode->paths[t2][i2].type);
|
||||
if (tx == GPU) srcNode->paths[t2][i2].type = PATH_PXN;
|
||||
srcNode->paths[t2][i2].bw = std::min(srcNode->paths[tx][ix].bw, cpuNode->paths[t2][i2].bw);
|
||||
return ncclSuccess;
|
||||
@@ -740,9 +747,9 @@ ncclResult_t ncclTopoComputePaths(struct ncclTopoSystem* system, struct ncclComm
|
||||
int c;
|
||||
NCCLCHECK(ncclGetLocalCpu(system, g, &c));
|
||||
if (c == -1) continue;
|
||||
if (gpuNode->paths[NET][n].type == PATH_PHB && gpuNode->paths[CPU][c].type == PATH_C2C) {
|
||||
gpuNode->paths[NET][n].type = PATH_P2C;
|
||||
netNode->paths[GPU][g].type = PATH_P2C;
|
||||
if (mergePathType(gpuNode->paths[CPU][c].type, netNode->paths[CPU][c].type) == PATH_P2C) {
|
||||
gpuNode->paths[NET][n].type = std::min(PATH_P2C, gpuNode->paths[NET][n].type);
|
||||
netNode->paths[GPU][g].type = std::min(PATH_P2C, netNode->paths[GPU][g].type);
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -781,16 +788,15 @@ ncclResult_t ncclTopoComputePaths(struct ncclTopoSystem* system, struct ncclComm
|
||||
// PXN = PCI + NVLink.
|
||||
struct ncclTopoNode* peerNode = system->nodes[GPU].nodes+localGpuIndex;
|
||||
// Only use PXN for NIC n if remote GPU p ...
|
||||
if (/* (1) is either connected to the NIC with PXB*/
|
||||
(peerNode->paths[NET][n].type <= PATH_PXB ||
|
||||
/* or with P2C and PxN over C2C is enabled */
|
||||
(ncclParamPxnC2c() && peerNode->paths[NET][n].type == PATH_P2C)) &&
|
||||
int pxnType = ncclParamPxnC2c() ? PATH_P2C : PATH_PXB;
|
||||
if (/* (1) is connected to the NIC with PxN type*/
|
||||
peerNode->paths[NET][n].type <= pxnType &&
|
||||
/* and (2) is connected to us through NVLink */
|
||||
peerNode->paths[GPU][g].type <= PATH_NVL &&
|
||||
/* and (3) is on the same node as us */
|
||||
NCCL_TOPO_ID_SYSTEM_ID(peerNode->id) == NCCL_TOPO_ID_SYSTEM_ID(gpu->id) &&
|
||||
/* and (4) has either higher bw to that NIC or avoid going through the CPU*/
|
||||
(peerNode->paths[NET][n].bw > gpu->paths[NET][n].bw || gpu->paths[NET][n].type > PATH_PXB))
|
||||
/* and (4) has either higher bw to that NIC or avoid going through the CPU (path.type is > PATH_PXN)*/
|
||||
(peerNode->paths[NET][n].bw > gpu->paths[NET][n].bw || gpu->paths[NET][n].type > PATH_PXN))
|
||||
// We can use that GPU as relay to communicate with that NIC.
|
||||
// Only enabling it in the GPU->NIC direction for now to favor
|
||||
// receiving locally and sending remotely (consistent with net.cc)
|
||||
@@ -811,6 +817,12 @@ ncclResult_t ncclTopoComputePaths(struct ncclTopoSystem* system, struct ncclComm
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// Pre-compute NET local gpus to accelerate search
|
||||
for (int n=0; n<system->nodes[NET].count; n++) {
|
||||
struct ncclTopoNode* net = system->nodes[NET].nodes+n;
|
||||
NCCLCHECK(ncclTopoGetLocalGpu(system, net->id, &net->net.localGpu));
|
||||
}
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
|
||||
@@ -474,6 +474,65 @@ ncclResult_t ncclTopoCompareGraphs(struct ncclTopoSystem* system, struct ncclTop
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
// Add the preferred NICs ordered by GPU first
|
||||
static ncclResult_t ncclTopoPrefNetsGpuFirst(struct ncclTopoSystem* system, int gpu, int nets[NCCL_TOPO_MAX_NODES], int* netCount) {
|
||||
const int nGpus = (gpu == -1) ? system->nodes[GPU].count : 1;
|
||||
int gpuCount = nGpus;
|
||||
int gpuIds[NCCL_TOPO_MAX_NODES] = {gpu};
|
||||
int firstNets[NCCL_TOPO_MAX_NODES];
|
||||
if (gpu == -1)
|
||||
for (int g = 0; g < nGpus; g++) gpuIds[g] = g;
|
||||
|
||||
for (int c = 0; c < MAXCHANNELS; c++) {
|
||||
for (int g = 0; g < nGpus; g++) {
|
||||
if (gpuIds[g] == -1) continue;
|
||||
int localNet;
|
||||
int64_t netId;
|
||||
struct ncclTopoNode* gpu = system->nodes[GPU].nodes + gpuIds[g];
|
||||
NCCLCHECK(ncclTopoGetLocalNet(system, gpu->gpu.rank, c, &netId, NULL));
|
||||
NCCLCHECK(ncclTopoIdToIndex(system, NET, netId, &localNet));
|
||||
// store the first net found for each GPU in case of duplicates
|
||||
if(c == 0) firstNets[g] = localNet;
|
||||
// if the NET has already been returned for channel 0, that GPU is done
|
||||
if (c > 0 && firstNets[g] == localNet) {
|
||||
gpuIds[g] = -1;
|
||||
gpuCount--;
|
||||
continue;
|
||||
}
|
||||
// only add it to the list if it doesn't already exist
|
||||
int found = 0;
|
||||
while (found < (*netCount) && nets[found] != localNet) found++;
|
||||
if (found == (*netCount)) nets[(*netCount)++] = localNet;
|
||||
}
|
||||
if (gpuCount == 0) break;
|
||||
}
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
// Add the preferred NICs ordered by channels first
|
||||
static ncclResult_t ncclTopoPrefNetsChannelFirst(struct ncclTopoSystem* system, int gpu, int nets[NCCL_TOPO_MAX_NODES], int* netCount) {
|
||||
for (int g = 0; g < system->nodes[GPU].count; g++) {
|
||||
if (gpu != -1 && gpu != g) continue;
|
||||
int localNetCount = 0, localNets[MAXCHANNELS];
|
||||
struct ncclTopoNode* gpu = system->nodes[GPU].nodes + g;
|
||||
for (int c = 0; c < MAXCHANNELS; c++) {
|
||||
int64_t netId;
|
||||
NCCLCHECK(ncclTopoGetLocalNet(system, gpu->gpu.rank, c, &netId, NULL));
|
||||
NCCLCHECK(ncclTopoIdToIndex(system, NET, netId, localNets + localNetCount));
|
||||
if (localNetCount > 0 && localNets[localNetCount] == localNets[0]) break;
|
||||
localNetCount++;
|
||||
}
|
||||
// Append NICs to list
|
||||
for (int i = 0; i < localNetCount; i++) {
|
||||
int n = localNets[i];
|
||||
int found = 0;
|
||||
while (found < (*netCount) && nets[found] != n) found++;
|
||||
if (found == (*netCount)) nets[(*netCount)++] = n;
|
||||
}
|
||||
}
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
// Build a sorted list of the NETs to try.
|
||||
//
|
||||
// "gpu" can be set to -1 to build a list suitable for all GPUs (search start) or to a given gpu
|
||||
@@ -482,39 +541,25 @@ ncclResult_t ncclTopoCompareGraphs(struct ncclTopoSystem* system, struct ncclTop
|
||||
// The list is built the following way:
|
||||
// 1. Select NETs starting with those close to GPU(s), based on paths[n].type.
|
||||
// 2. add other NETs satisfying typeInter but not already in the list.
|
||||
|
||||
NCCL_PARAM(ScatterEnable, "MNNVL_SCATTER_NETS_ENABLE", 1);
|
||||
ncclResult_t ncclTopoSelectNets(struct ncclTopoSystem* system, int typeInter, int gpu, int nets[NCCL_TOPO_MAX_NODES], int* netCountRet) {
|
||||
ncclResult_t ret = ncclSuccess;
|
||||
int netCount = 0;
|
||||
int localNetCount;
|
||||
int localNets[MAXCHANNELS];
|
||||
|
||||
// First add the preferred NICs
|
||||
for (int g=0; g<system->nodes[GPU].count; g++) {
|
||||
if (gpu != -1 && gpu != g) continue;
|
||||
localNetCount = 0;
|
||||
struct ncclTopoNode* gpu = system->nodes[GPU].nodes+g;
|
||||
for (int c = 0; c<MAXCHANNELS; c++) {
|
||||
int64_t netId;
|
||||
NCCLCHECK(ncclTopoGetLocalNet(system, gpu->gpu.rank, c, &netId, NULL));
|
||||
NCCLCHECK(ncclTopoIdToIndex(system, NET, netId, localNets+localNetCount));
|
||||
if (localNetCount > 0 && localNets[localNetCount] == localNets[0]) break;
|
||||
localNetCount++;
|
||||
}
|
||||
// Append NICs to list
|
||||
for (int i=0; i<localNetCount; i++) {
|
||||
int n = localNets[i];
|
||||
int found = 0;
|
||||
while (found<netCount && nets[found] != n) found++;
|
||||
if (found == netCount) nets[netCount++] = n;
|
||||
}
|
||||
// First add the preferred NETs.
|
||||
if (system->nHosts > 1 && ncclParamScatterEnable()) {
|
||||
// For MNNVL systems, we sort the devices by GPU first, then by channel
|
||||
NCCLCHECK(ncclTopoPrefNetsGpuFirst(system, gpu, nets, &netCount));
|
||||
} else {
|
||||
// For other systems, we sort the devices by channel first, then by GPU
|
||||
NCCLCHECK(ncclTopoPrefNetsChannelFirst(system, gpu, nets, &netCount));
|
||||
}
|
||||
|
||||
// Then add others satisfying typeInter
|
||||
for (int t=0; t <= typeInter; t++) {
|
||||
for (int g=0; g<system->nodes[GPU].count; g++) {
|
||||
for (int g = 0; g < system->nodes[GPU].count; g++) {
|
||||
if (gpu != -1 && gpu != g) continue;
|
||||
localNetCount = 0;
|
||||
int localNetCount = 0, localNets[MAXCHANNELS];
|
||||
struct ncclTopoNode* gpu = system->nodes[GPU].nodes+g;
|
||||
struct ncclTopoLinkList* paths = gpu->paths[NET];
|
||||
for (int n=0; n<system->nodes[NET].count && n<MAXCHANNELS; n++) {
|
||||
@@ -662,8 +707,7 @@ ncclResult_t ncclTopoSearchRecNet(struct ncclTopoSystem* system, struct ncclTopo
|
||||
if (graph->pattern == NCCL_TOPO_PATTERN_NVLS || graph->pattern == NCCL_TOPO_PATTERN_COLLNET_DIRECT) {
|
||||
// NVLS search only tries to find NIC:GPU combinations to compute the heads.
|
||||
if (graph->nChannels < netCount) {
|
||||
int gpu;
|
||||
NCCLCHECK(ncclTopoGetLocalGpu(system, net->id, &gpu));
|
||||
int gpu = net->net.localGpu;
|
||||
if (gpu != -1) {
|
||||
int duplicate = 0;
|
||||
// check whether there is duplicate head when one GPU connects with multiple NICs
|
||||
@@ -680,13 +724,12 @@ ncclResult_t ncclTopoSearchRecNet(struct ncclTopoSystem* system, struct ncclTopo
|
||||
}
|
||||
}
|
||||
} else {
|
||||
if (graph->nChannels > 0) {
|
||||
if (graph->nChannels > 0 && graph->sameChannels == 1) {
|
||||
// Try to replay the last channel
|
||||
int g;
|
||||
NCCLCHECK(ncclTopoReplayGetGpu(system, graph, -1, &g));
|
||||
NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, 0, backToNet, backToFirstRank, FORCED_ORDER_REPLAY, time, NET, n, g));
|
||||
}
|
||||
if (graph->nChannels == 0 || graph->sameChannels == 0) {
|
||||
} else {
|
||||
if (graph->nChannels == 0 && system->nodes[NVS].count == 0) {
|
||||
// Always try the PCI order first to set a reference, but don't count in the timeout nor let it run for long
|
||||
struct ncclTopoLinkList* paths = net->paths[GPU];
|
||||
@@ -709,11 +752,16 @@ ncclResult_t ncclTopoSearchRecNet(struct ncclTopoSystem* system, struct ncclTopo
|
||||
}
|
||||
|
||||
// Then try the most local GPUs
|
||||
int localGpu = net->net.localGpu;
|
||||
if (localGpu != -1) {
|
||||
NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, 0, backToNet, backToFirstRank, 0, time, NET, n, localGpu));
|
||||
}
|
||||
int localGpus[NCCL_TOPO_MAX_NODES], localGpuCount, pathType;
|
||||
NCCLCHECK(ncclTopoGetLocal(system, NET, n, GPU, localGpus, &localGpuCount, &pathType));
|
||||
// if no GPUs are connected, skip this net
|
||||
if (pathType == PATH_DIS) continue;
|
||||
for (int g = 0; g < localGpuCount; ++g) {
|
||||
if (localGpus[g] == localGpu) continue; // We already tried this one
|
||||
NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, 0, backToNet, backToFirstRank, 0, time, NET, n, localGpus[g]));
|
||||
}
|
||||
}
|
||||
@@ -800,8 +848,8 @@ struct kvDict kvDictLinkType[] = {
|
||||
{ "NVB", PATH_NVB },
|
||||
{ "PIX", PATH_PIX },
|
||||
{ "PXB", PATH_PXB },
|
||||
{ "PXN", PATH_PXN },
|
||||
{ "P2C", PATH_P2C },
|
||||
{ "PXN", PATH_PXN },
|
||||
{ "PHB", PATH_PHB },
|
||||
{ "SYS", PATH_SYS },
|
||||
{ NULL, 0 }
|
||||
@@ -849,8 +897,10 @@ ncclResult_t ncclTopoGetGraphFromXmlSub(struct ncclXmlNode *xmlGraph, struct ncc
|
||||
NCCLCHECK(xmlGetAttrInt(xmlGraph, "nchannels", &graph->nChannels));
|
||||
NCCLCHECK(xmlGetAttrFloat(xmlGraph, "speedintra", &graph->bwIntra));
|
||||
NCCLCHECK(xmlGetAttrFloat(xmlGraph, "speedinter", &graph->bwInter));
|
||||
if (xmlGetAttrFloat(xmlGraph, "latencyinter", &graph->latencyInter) != ncclSuccess) graph->latencyInter = 0.0;
|
||||
const char* str;
|
||||
NCCLCHECK(xmlGetAttr(xmlGraph, "latencyinter", &str));
|
||||
if (!str) INFO(NCCL_GRAPH, "latencyinter not found in graph, using 0.0");
|
||||
graph->latencyInter = str ? strtof(str, NULL) : 0.0;
|
||||
NCCLCHECK(xmlGetAttr(xmlGraph, "typeintra", &str));
|
||||
NCCLCHECK(kvConvertToInt(str, &graph->typeIntra, kvDictLinkType));
|
||||
NCCLCHECK(xmlGetAttr(xmlGraph, "typeinter", &str));
|
||||
@@ -970,7 +1020,7 @@ float sm90SpeedArrayInter[] = { 48.0, 45.0, 42.0, 40.0, 30.0, 24.0, 22.0, 20.0,
|
||||
RCCL_PARAM(ModelMatchingDisable, "MODEL_MATCHING_DISABLE", 0);
|
||||
|
||||
float sm100SpeedArrayIntra[] = { 90.0, 80.0, 70.0, 60.0, 50.0, 40.0, 30.0, 24.0, 20.0, 19.0, 18.0 };
|
||||
float sm100SpeedArrayInter[] = { 47.9, 45.0, 42.0, 40.0, 30.0, 24.0, 22.0, 20.0, 17.5, 15.0, 12.0, 6.0, 3.0, 2.4, 1.2, 0.24, 0.12 };
|
||||
float sm100SpeedArrayInter[] = { 96.0, 48.0, 45.1, 42.0, 40.0, 30.0, 24.0, 22.0, 20.0, 17.5, 15.0, 12.0, 6.0, 3.0, 2.4, 1.2, 0.24, 0.12 };
|
||||
#define NSPEEDSINTRA_SM100 (sizeof(sm100SpeedArrayIntra)/sizeof(float))
|
||||
#define NSPEEDSINTER_SM100 (sizeof(sm100SpeedArrayInter)/sizeof(float))
|
||||
|
||||
@@ -1249,13 +1299,17 @@ ncclResult_t ncclTopoPrintGraph(struct ncclTopoSystem* system, struct ncclTopoGr
|
||||
sprintf(line+offset, " NET/%d", n);
|
||||
offset = strlen(line);
|
||||
}
|
||||
sprintf(line+offset, " %s/%d", topoNodeTypeStr[GPU], graph->intra[ngpus*c+i]);
|
||||
int g;
|
||||
ncclTopoRankToIndex(system, graph->intra[ngpus * c + i], &g, true);
|
||||
int64_t topoId = system->nodes[GPU].nodes[g].id;
|
||||
sprintf(line + offset, " %s/%lx-%lx", topoNodeTypeStr[GPU], NCCL_TOPO_ID_SYSTEM_ID(topoId), NCCL_TOPO_ID_LOCAL_ID(topoId));
|
||||
offset = strlen(line);
|
||||
n = graph->intraNets[(ngpus*c+i)*2+1]-'N';
|
||||
if(n >= 0 && n < system->nodes[NET].count) {
|
||||
sprintf(line+offset, " NET/%d", n);
|
||||
offset = strlen(line);
|
||||
}
|
||||
if (graph->id == 3) break; // NVLS graphs only use the first GPU
|
||||
}
|
||||
if (system->nodes[NET].count > 0 && system->nodes[GPU].count != system->nRanks && !graph->nIntraChannels) {
|
||||
sprintf(line+offset, " %s/%lx-%lx", topoNodeTypeStr[NET], NCCL_TOPO_ID_SYSTEM_ID(graph->inter[2*c+1]), NCCL_TOPO_ID_LOCAL_ID(graph->inter[2*c+1]));
|
||||
@@ -1371,7 +1425,8 @@ ncclResult_t ncclTopoGetNetDev(struct ncclComm* comm, int rank, struct ncclTopoG
|
||||
NCCLCHECK(ncclTopoGetLocalGpu(comm->topo, netId, &g2));
|
||||
if (g2 != -1) {
|
||||
struct ncclTopoNode* peerGpu = comm->topo->nodes[GPU].nodes+g2;
|
||||
if (peerGpu->paths[GPU][g1].type <= PATH_NVL && peerGpu->paths[NET][n].type <= PATH_PXB) {
|
||||
int pxnType = ncclParamPxnC2c() ? PATH_P2C : PATH_PXB;
|
||||
if (peerGpu->paths[GPU][g1].type <= PATH_NVL && peerGpu->paths[NET][n].type <= pxnType) {
|
||||
*proxyRank = peerGpu->gpu.rank;
|
||||
if (dev) *dev = netDev;
|
||||
if (id) *id = netId;
|
||||
|
||||
@@ -23,10 +23,10 @@
|
||||
const char* topoNodeTypeStr[] = { "GPU", "PCI", "NVS", "CPU", "NIC", "NET" };
|
||||
#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
|
||||
const char* topoLinkTypeStr[] = { "LOC", "XGMI", "", "C2C", "PCI", "", "", "", "", "SYS", "NET" };
|
||||
const char* topoPathTypeStr[] = { "LOC", "XGMI", "NVB", "C2C", "PIX", "PXB", "PXN", "P2C", "PHB", "SYS", "NET", "DIS" };
|
||||
const char* topoPathTypeStr[] = { "LOC", "XGMI", "NVB", "C2C", "PIX", "PXB", "P2C", "PXN", "PHB", "SYS", "NET", "DIS" };
|
||||
#else
|
||||
const char* topoLinkTypeStr[] = { "LOC", "NVL", "", "C2C", "PCI", "", "", "", "", "SYS", "NET" };
|
||||
const char* topoPathTypeStr[] = { "LOC", "NVL", "NVB", "C2C", "PIX", "PXB", "PXN", "P2C", "PHB", "SYS", "NET", "DIS" };
|
||||
const char* topoPathTypeStr[] = { "LOC", "NVL", "NVB", "C2C", "PIX", "PXB", "P2C", "PXN", "PHB", "SYS", "NET", "DIS" };
|
||||
#endif
|
||||
|
||||
/******************************************************************/
|
||||
@@ -760,7 +760,14 @@ ncclResult_t ncclTopoGetSystemFromXml(struct ncclXml* xml, struct ncclTopoSystem
|
||||
struct ncclXmlNode* node = topNode->subs[s];
|
||||
if (strcmp(node->name, "cpu") == 0) NCCLCHECK(ncclTopoAddCpu(node, *topoSystem));
|
||||
}
|
||||
for (int systemId=0; systemId<system->nHosts; systemId++) if (system->hostHashes[systemId] == localHostHash) system->systemId = systemId;
|
||||
|
||||
int systemId = 0;
|
||||
while (systemId < system->nHosts && system->hostHashes[systemId] != localHostHash) systemId++;
|
||||
system->systemId = systemId;
|
||||
if(systemId == system->nHosts){
|
||||
WARN("localHostHash = 0x%lx not found in the list of system hostHashes",localHostHash);
|
||||
return ncclInvalidArgument;
|
||||
}
|
||||
|
||||
#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
|
||||
NCCLCHECK(ncclTopoAddXGMI(topNode, *topoSystem, NULL));
|
||||
@@ -1230,8 +1237,8 @@ struct kvDict nicPathKvList[] = {
|
||||
{ "PORT", PATH_PORT },
|
||||
{ "PIX", PATH_PIX },
|
||||
{ "PXB", PATH_PXB },
|
||||
{ "PXN", PATH_PXN },
|
||||
{ "P2C", PATH_P2C },
|
||||
{ "PXN", PATH_PXN },
|
||||
{ "PHB", PATH_PHB },
|
||||
{ "SYS", PATH_SYS },
|
||||
{ NULL, 0 }
|
||||
@@ -1509,7 +1516,7 @@ ncclResult_t ncclTopoGetSystem(struct ncclComm* comm, struct ncclTopoSystem** sy
|
||||
}
|
||||
|
||||
// Only update our topo tracking structure if we aren't dumping (separate steps)
|
||||
if (dumpXmlFile == NULL) NCCLCHECKGOTO(ncclTopoGetSystemFromXml(xml, system, comm->peerInfo[comm->rank].hostHash), ret, fail);
|
||||
if (dumpXmlFile == NULL) NCCLCHECKGOTO(ncclTopoGetSystemFromXml(xml, system, getHostHash()), ret, fail);
|
||||
|
||||
exit:
|
||||
if (!comm->MNNVL && localRanks) free(localRanks);
|
||||
|
||||
@@ -21,7 +21,7 @@
|
||||
#define SM80_NVLINK_BW 20.0
|
||||
#define SM90_NVLINK_BW 20.6
|
||||
#define SM86_NVLINK_BW 12.0
|
||||
#define SM100_NVLINK_BW 40.0
|
||||
#define SM100_NVLINK_BW 40.1
|
||||
#define PCI_BW 12.0 // PCI Gen3 x16
|
||||
#define AMD_BW 16.0
|
||||
#define BDW_QPI_BW 6.0
|
||||
@@ -83,11 +83,11 @@ extern const char* topoLinkTypeStr[];
|
||||
// Connection traversing multiple PCIe bridges (without traversing the PCIe Host Bridge)
|
||||
#define PATH_PXB 5
|
||||
|
||||
// Connection between a GPU and a NIC using an intermediate GPU. Used to enable rail-local, aggregated network send/recv operations.
|
||||
#define PATH_PXN 6
|
||||
|
||||
// Connection between a GPU and a NIC using the C2C connection to the CPU and the PCIe connection to the NIC
|
||||
#define PATH_P2C 7
|
||||
#define PATH_P2C 6
|
||||
|
||||
// Connection between a GPU and a NIC using an intermediate GPU. Used to enable rail-local, aggregated network send/recv operations.
|
||||
#define PATH_PXN 7
|
||||
|
||||
// Connection traversing PCIe as well as a PCIe Host Bridge (typically the CPU)
|
||||
#define PATH_PHB 8
|
||||
@@ -105,6 +105,8 @@ extern const char* topoLinkTypeStr[];
|
||||
#define PATH_DIS 11
|
||||
extern const char* topoPathTypeStr[];
|
||||
|
||||
extern int64_t ncclParamPxnC2c();
|
||||
|
||||
struct ncclTopoNode;
|
||||
struct ncclTopoLink {
|
||||
int type;
|
||||
@@ -163,6 +165,7 @@ struct ncclTopoNode {
|
||||
int gdrSupport;
|
||||
int collSupport;
|
||||
int maxChannels;
|
||||
int localGpu;
|
||||
int64_t busId;
|
||||
}net;
|
||||
struct {
|
||||
|
||||
@@ -748,9 +748,16 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom
|
||||
pEnable = 0;
|
||||
#endif
|
||||
#else
|
||||
// Enable LL128 by default only on Volta/Ampere/Hopper+NVLink. Other cases are not tested and may cause silent data corruption.
|
||||
pEnable = 1;
|
||||
pEnable &= (graphs[a]->typeInter <= PATH_PXB || (minCompCap >= 90 && graphs[a]->typeInter <= (ncclParamLl128C2c() ? PATH_P2C : PATH_PXN)));
|
||||
if (ncclParamLl128C2c() && minCompCap >= 90) {
|
||||
// Enable LL128 by default only on Hopper/Blackwell for all connections up to P2C and PXN.
|
||||
pEnable &= (graphs[a]->typeInter <= PATH_PXN);
|
||||
} else {
|
||||
// Enable LL128 only up to PXB. Don't enable LL128 over PxN because PxN can encapsulate PxB or P2C links.
|
||||
pEnable &= (graphs[a]->typeInter <= PATH_PXB);
|
||||
if (!ncclParamLl128C2c() && minCompCap >= 90)
|
||||
INFO(NCCL_GRAPH, "Disabling LL128 over all PxN connections (PXB and C2C). This ensures that no C2C link will be used by LL128.");
|
||||
}
|
||||
pEnable &= (graphs[a]->typeIntra <= PATH_NVB);
|
||||
pEnable &= (minCompCap == maxCompCap);
|
||||
pEnable &= !(minCompCap < 70 || (minCompCap == 90 && CUDART_VERSION == 11080 && c == ncclFuncAllReduce && a == NCCL_ALGO_RING && comm->nRanks == 2));
|
||||
|
||||
@@ -9,6 +9,7 @@
|
||||
#include <stdint.h>
|
||||
#include <sys/types.h>
|
||||
#include <unistd.h>
|
||||
#include <string.h>
|
||||
|
||||
#if __GNUC__ >= 3
|
||||
# define __attribute_const __attribute__((const))
|
||||
@@ -39,7 +40,7 @@ union ibv_gid {
|
||||
#define vext_field_avail(type, fld, sz) (offsetof(type, fld) < (sz))
|
||||
|
||||
/*XXX:__VERBS_ABI_IS_EXTENDED produces warning "integer operation result is out of range" with g++ 4.8.2*/
|
||||
//static void *__VERBS_ABI_IS_EXTENDED = ((uint8_t *)NULL) - 1;
|
||||
static void *__VERBS_ABI_IS_EXTENDED = ((uint8_t *)NULL) - 1;
|
||||
|
||||
enum ibv_node_type {
|
||||
IBV_NODE_UNKNOWN = -1,
|
||||
@@ -208,7 +209,9 @@ struct ibv_port_attr {
|
||||
uint8_t active_speed;
|
||||
uint8_t phys_state;
|
||||
uint8_t link_layer;
|
||||
uint8_t reserved;
|
||||
uint8_t flags;
|
||||
uint16_t port_cap_flags2;
|
||||
uint32_t active_speed_ex;
|
||||
};
|
||||
|
||||
enum ibv_event_type {
|
||||
@@ -993,37 +996,50 @@ enum verbs_context_mask {
|
||||
|
||||
struct verbs_context {
|
||||
/* "grows up" - new fields go here */
|
||||
int (*_reserved_2) (void);
|
||||
int (*destroy_flow) (struct ibv_flow *flow);
|
||||
int (*_reserved_1) (void);
|
||||
struct ibv_flow * (*create_flow) (struct ibv_qp *qp,
|
||||
struct ibv_flow_attr *flow_attr);
|
||||
int (*query_port)(struct ibv_context *context, uint8_t port_num,
|
||||
struct ibv_port_attr *port_attr,
|
||||
size_t port_attr_len);
|
||||
int (*_reserved[25]) (void);
|
||||
struct verbs_ex_private *priv;
|
||||
int (*query_device_ex)(struct ibv_context *context,
|
||||
const struct ibv_query_device_ex_input *input,
|
||||
struct ibv_device_attr_ex *attr,
|
||||
size_t attr_size);
|
||||
int (*ibv_destroy_flow) (struct ibv_flow *flow);
|
||||
void (*ABI_placeholder2) (void); /* DO NOT COPY THIS GARBAGE */
|
||||
struct ibv_flow * (*ibv_create_flow) (struct ibv_qp *qp,
|
||||
struct ibv_flow_attr *flow_attr);
|
||||
void (*ABI_placeholder1) (void); /* DO NOT COPY THIS GARBAGE */
|
||||
struct ibv_qp * (*open_qp)(struct ibv_context *context,
|
||||
struct ibv_qp_open_attr *attr);
|
||||
struct ibv_qp * (*create_qp_ex)(struct ibv_context *context,
|
||||
struct ibv_qp_init_attr_ex *qp_init_attr_ex);
|
||||
int (*get_srq_num)(struct ibv_srq *srq, uint32_t *srq_num);
|
||||
struct ibv_srq * (*create_srq_ex)(struct ibv_context *context,
|
||||
struct ibv_srq_init_attr_ex *srq_init_attr_ex);
|
||||
struct ibv_xrcd * (*open_xrcd)(struct ibv_context *context,
|
||||
struct ibv_xrcd_init_attr *xrcd_init_attr);
|
||||
int (*close_xrcd)(struct ibv_xrcd *xrcd);
|
||||
uint64_t has_comp_mask;
|
||||
size_t sz; /* Must be immediately before struct ibv_context */
|
||||
struct ibv_context context;/* Must be last field in the struct */
|
||||
struct ibv_srq * (*create_srq_ex)(struct ibv_context *context,
|
||||
struct ibv_srq_init_attr_ex *srq_init_attr_ex);
|
||||
struct ibv_xrcd * (*open_xrcd)(struct ibv_context *context,
|
||||
struct ibv_xrcd_init_attr *xrcd_init_attr);
|
||||
int (*close_xrcd)(struct ibv_xrcd *xrcd);
|
||||
uint64_t _ABI_placeholder3;
|
||||
size_t sz; /* Must be immediately before struct ibv_context */
|
||||
struct ibv_context context; /* Must be last field in the struct */
|
||||
};
|
||||
|
||||
/*XXX:__VERBS_ABI_IS_EXTENDED produces warning "integer operation result is out of range" with g++ 4.8.2*/
|
||||
/*static inline struct verbs_context *verbs_get_ctx(struct ibv_context *ctx)
|
||||
static inline struct verbs_context *verbs_get_ctx(struct ibv_context *ctx)
|
||||
{
|
||||
return (!ctx || (ctx->abi_compat != __VERBS_ABI_IS_EXTENDED)) ?
|
||||
NULL : container_of(ctx, struct verbs_context, context);
|
||||
if (ctx->abi_compat != __VERBS_ABI_IS_EXTENDED)
|
||||
return NULL;
|
||||
|
||||
/* open code container_of to not pollute the global namespace */
|
||||
return (struct verbs_context *)(((uintptr_t)ctx) -
|
||||
offsetof(struct verbs_context,
|
||||
context));
|
||||
}
|
||||
|
||||
#define verbs_get_ctx_op(ctx, op) ({ \
|
||||
struct verbs_context *_vctx = verbs_get_ctx(ctx); \
|
||||
(!_vctx || (_vctx->sz < sizeof(*_vctx) - offsetof(struct verbs_context, op)) || \
|
||||
!_vctx->op) ? NULL : _vctx; })*/
|
||||
struct verbs_context *__vctx = verbs_get_ctx(ctx); \
|
||||
(!__vctx || (__vctx->sz < sizeof(*__vctx) - offsetof(struct verbs_context, op)) || \
|
||||
!__vctx->op) ? NULL : __vctx; })
|
||||
|
||||
#define verbs_set_ctx_op(_vctx, op, ptr) ({ \
|
||||
struct verbs_context *vctx = _vctx; \
|
||||
@@ -1055,4 +1071,20 @@ struct ibv_ece {
|
||||
uint32_t comp_mask;
|
||||
};
|
||||
|
||||
/**
|
||||
* ibv_query_port_ex - Get (extended) port properties
|
||||
*/
|
||||
static inline int ibv_query_port_ex(struct ibv_context *context,
|
||||
uint8_t port_num,
|
||||
struct ibv_port_attr *port_attr)
|
||||
{
|
||||
struct verbs_context *vctx = verbs_get_ctx_op(context, query_port);
|
||||
|
||||
if (vctx) {
|
||||
return vctx->query_port(context, port_num, port_attr, sizeof(*port_attr));
|
||||
}
|
||||
|
||||
return -1;
|
||||
}
|
||||
|
||||
#endif // NCCL_IBV_CORE_H_
|
||||
|
||||
@@ -9,10 +9,16 @@
|
||||
|
||||
#include "nccl.h"
|
||||
|
||||
enum ncclPluginType {
|
||||
ncclPluginTypeNet,
|
||||
ncclPluginTypeTuner,
|
||||
ncclPluginTypeProfiler,
|
||||
};
|
||||
|
||||
void* ncclOpenNetPluginLib(const char* name);
|
||||
void* ncclOpenTunerPluginLib(const char* name);
|
||||
void* ncclOpenProfilerPluginLib(const char* name);
|
||||
void* ncclGetNetPluginLib(void);
|
||||
ncclResult_t ncclClosePluginLib(void* handle);
|
||||
void* ncclGetNetPluginLib(enum ncclPluginType type);
|
||||
ncclResult_t ncclClosePluginLib(void* handle, enum ncclPluginType type);
|
||||
|
||||
#endif
|
||||
|
||||
@@ -2154,7 +2154,7 @@ static ncclResult_t envConfigOverride(ncclComm_t comm) {
|
||||
int minCTAsEnv;
|
||||
int maxCTAsEnv;
|
||||
int splitShareEnv;
|
||||
int collnetEnableEnv;
|
||||
const char* collnetEnableEnv;
|
||||
int ctaPolicyEnv;
|
||||
int shrinkShareEnv;
|
||||
int nvlsCTAsEnv;
|
||||
@@ -2208,9 +2208,15 @@ static ncclResult_t envConfigOverride(ncclComm_t comm) {
|
||||
comm->config.shrinkShare = shrinkShareEnv;
|
||||
}
|
||||
|
||||
collnetEnableEnv = ncclParamCollnetEnable();
|
||||
if (collnetEnableEnv != NCCL_CONFIG_UNDEF_INT) {
|
||||
comm->config.collnetEnable = collnetEnableEnv;
|
||||
// NCCL_COLLNET_ENABLE needs to be reloaded each time for comm init
|
||||
// since users might change the env on the fly to enable/disable collnet
|
||||
collnetEnableEnv = ncclGetEnv("NCCL_COLLNET_ENABLE");
|
||||
if (collnetEnableEnv != NULL) {
|
||||
int collnetEnableInt = (int)strtol(collnetEnableEnv, NULL, 0);
|
||||
if (collnetEnableInt != NCCL_CONFIG_UNDEF_INT) {
|
||||
comm->config.collnetEnable = collnetEnableInt;
|
||||
INFO(NCCL_ENV, "NCCL_COLLNET_ENABLE set by environment to %d.", collnetEnableInt);
|
||||
}
|
||||
}
|
||||
|
||||
ctaPolicyEnv = ncclParamCtaPolicy();
|
||||
@@ -2867,6 +2873,7 @@ ncclResult_t ncclCommDestroy_impl(ncclComm_t comm) {
|
||||
NVTX3_PAYLOAD(comm->commHash, nranks, rank, cudaDev));
|
||||
|
||||
TRACE(NCCL_INIT, "comm %p rank %d nRanks %d cudaDev %d busId %lx", comm, rank, nranks, cudaDev, comm->busId);
|
||||
NCCLCHECK(ncclGroupStartInternal());
|
||||
// Try and prevent a double free of the comm struct (user error)
|
||||
if (comm->rank == -1 || comm->nRanks == -1 || comm->cudaDev == -1 || comm->busId == -1) {
|
||||
WARN("comm %p has already been destroyed", comm);
|
||||
@@ -2881,6 +2888,8 @@ ncclResult_t ncclCommDestroy_impl(ncclComm_t comm) {
|
||||
NCCLCHECKGOTO(ncclAsyncLaunch((struct ncclAsyncJob*)job, commReclaim, NULL, free, comm), res, fail);
|
||||
|
||||
exit:
|
||||
ncclGroupErrCheck(res);
|
||||
NCCLCHECK(ncclGroupEndInternal());
|
||||
return res;
|
||||
fail:
|
||||
goto exit;
|
||||
@@ -2905,6 +2914,7 @@ ncclResult_t ncclCommAbort_impl(ncclComm_t comm) {
|
||||
if (comm == NULL) {
|
||||
return ncclSuccess;
|
||||
}
|
||||
NCCLCHECK(ncclGroupStartInternal());
|
||||
// Ask anything that might still be running on the device to quit
|
||||
NCCLCHECK(setCommAbortFlags(comm,1));
|
||||
comm->destroyFlag = 1;
|
||||
@@ -2927,7 +2937,9 @@ ncclResult_t ncclCommAbort_impl(ncclComm_t comm) {
|
||||
NCCLCHECKGOTO(ncclAsyncLaunch((struct ncclAsyncJob*)job, commReclaim, NULL, free, comm), res, fail);
|
||||
|
||||
exit:
|
||||
return ncclSuccess;
|
||||
ncclGroupErrCheck(res);
|
||||
NCCLCHECK(ncclGroupEndInternal());
|
||||
return res;
|
||||
fail:
|
||||
goto exit;
|
||||
}
|
||||
|
||||
@@ -143,8 +143,14 @@ ncclResult_t wrap_ibv_query_device(struct ibv_context *context, struct ibv_devic
|
||||
IBV_INT_CHECK_RET_ERRNO(ibvSymbols, ibv_internal_query_device, ibv_internal_query_device(context, device_attr), 0, "ibv_query_device");
|
||||
}
|
||||
|
||||
ncclResult_t wrap_ibv_query_port(struct ibv_context *context, uint8_t port_num, struct ibv_port_attr *port_attr) { /*returns 0 on success, or the value of errno on failure (which indicates the failure reason)*/
|
||||
IBV_INT_CHECK_RET_ERRNO(ibvSymbols, ibv_internal_query_port, ibv_internal_query_port(context, port_num, port_attr), 0, "ibv_query_port");
|
||||
ncclResult_t wrap_ibv_query_port(struct ibv_context *context, uint8_t port_num, struct ibv_port_attr *port_attr) {
|
||||
// First try and query the extended port attributes (e.g. active_speed_ex)
|
||||
if (ibv_query_port_ex(context, port_num, port_attr) != 0) {
|
||||
// Fall back to the original attribute API call, but zero all members first
|
||||
memset(port_attr, 0, sizeof(*port_attr));
|
||||
IBV_INT_CHECK_RET_ERRNO(ibvSymbols, ibv_internal_query_port, ibv_internal_query_port(context, port_num, port_attr), 0, "ibv_query_port");
|
||||
}
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
ncclResult_t wrap_ibv_query_gid(struct ibv_context *context, uint8_t port_num, int index, union ibv_gid *gid) {
|
||||
|
||||
@@ -52,6 +52,9 @@ ncclResult_t buildMlx5dvSymbols(struct ncclMlx5dvSymbols* mlx5dvSymbols) {
|
||||
#define LOAD_SYM_VERSION(handle, symbol, funcptr, version) do { \
|
||||
cast = (void**)&funcptr; \
|
||||
*cast = dlvsym(handle, symbol, version); \
|
||||
if (*cast == NULL) { \
|
||||
INFO(NCCL_NET, "dlvsym failed on %s - %s version %s", symbol, dlerror(), version); \
|
||||
} \
|
||||
} while (0)
|
||||
|
||||
LOAD_SYM(mlx5dvhandle, "mlx5dv_is_supported", mlx5dvSymbols->mlx5dv_internal_is_supported);
|
||||
|
||||
@@ -446,7 +446,8 @@ static ncclResult_t socketTryAccept(struct ncclSocket* sock) {
|
||||
if (sock->fd != -1) {
|
||||
sock->state = ncclSocketStateAccepted;
|
||||
} else if (errno == ENETDOWN || errno == EPROTO || errno == ENOPROTOOPT || errno == EHOSTDOWN ||
|
||||
errno == ENONET || errno == EHOSTUNREACH || errno == EOPNOTSUPP || errno == ENETUNREACH) {
|
||||
errno == ENONET || errno == EHOSTUNREACH || errno == EOPNOTSUPP || errno == ENETUNREACH ||
|
||||
errno == EINTR) {
|
||||
/* per accept's man page, for linux sockets, the following errors might be already pending errors
|
||||
* and should be considered as EAGAIN. To avoid infinite loop in case of errors, we use the retry count*/
|
||||
if (++sock->errorRetries == ncclParamRetryCnt()) {
|
||||
|
||||
@@ -21,7 +21,6 @@ struct ncclStrongStreamCapture {
|
||||
cudaGraph_t graph;
|
||||
unsigned long long graphId;
|
||||
cudaStream_t captureStream;
|
||||
cudaGraphNode_t lastRecord;
|
||||
void* acquiredBy;
|
||||
};
|
||||
|
||||
@@ -194,7 +193,6 @@ ncclResult_t ncclStrongStreamAcquire(
|
||||
CUDACHECKGOTO(cudaStreamCreateWithFlags(&cap->captureStream, cudaStreamNonBlocking), ret, do_unlock);
|
||||
}
|
||||
cap->graphId = graph.graphId;
|
||||
cap->lastRecord = nullptr;
|
||||
cap->acquiredBy = localThreadId();
|
||||
// Push to capturing list.
|
||||
cap->next = ss->captureHead;
|
||||
@@ -274,16 +272,6 @@ ncclResult_t ncclStrongStreamRelease(
|
||||
cudaGraphNode_t recordNode;
|
||||
CUDACHECK(cudaGraphAddEventRecordNode(&recordNode, graph.graph, nullptr, 0, ss->serialEvent));
|
||||
|
||||
// Make this record order after previous record on this stream.
|
||||
if (cap->lastRecord != nullptr) {
|
||||
#if CUDART_VERSION >= 13000
|
||||
CUDACHECK(cudaGraphAddDependencies_v2(graph.graph, &cap->lastRecord, &recordNode, nullptr, 1));
|
||||
#else
|
||||
CUDACHECK(cudaGraphAddDependencies(graph.graph, &cap->lastRecord, &recordNode, 1));
|
||||
#endif
|
||||
}
|
||||
cap->lastRecord = recordNode;
|
||||
|
||||
// Get current nodes from work stream so we can add them as dependencies.
|
||||
cudaStreamCaptureStatus status;
|
||||
cudaGraphNode_t const* nodes;
|
||||
@@ -316,6 +304,22 @@ ncclResult_t ncclStrongStreamRelease(
|
||||
}
|
||||
}
|
||||
|
||||
// Make every future operation captured on cap->captureStream depend on 'recordNode'.
|
||||
#if CUDART_VERSION >= 13000
|
||||
CUDACHECK(cudaStreamUpdateCaptureDependencies_v2(
|
||||
cap->captureStream,
|
||||
&recordNode, /* dependencies */
|
||||
/*edges =*/ nullptr, /* no edge annotations */
|
||||
1, /* count */
|
||||
cudaStreamSetCaptureDependencies));
|
||||
#else
|
||||
CUDACHECK(cudaStreamUpdateCaptureDependencies(
|
||||
cap->captureStream,
|
||||
&recordNode,
|
||||
1,
|
||||
cudaStreamSetCaptureDependencies));
|
||||
#endif
|
||||
|
||||
if (cap->acquiredBy != localThreadId() && ncclParamLaunchRaceFatal()) {
|
||||
WARN("%s", launchRaceFatalMsg);
|
||||
return ncclInvalidUsage;
|
||||
|
||||
@@ -67,7 +67,7 @@ static pthread_once_t initPluginLibsOnceControl = PTHREAD_ONCE_INIT;
|
||||
static ncclResult_t ncclNetPluginUnload(netPluginLib_t* pluginLib) {
|
||||
if ((pluginLib->dlHandle) && ((pluginLib->ncclNetPluginRefCount) == 0)) {
|
||||
INFO(NCCL_INIT|NCCL_NET, "Unloading plugin %s", pluginLib->name);
|
||||
NCCLCHECK(ncclClosePluginLib(pluginLib->dlHandle));
|
||||
NCCLCHECK(ncclClosePluginLib(pluginLib->dlHandle, ncclPluginTypeNet));
|
||||
memset(pluginLib, 0, sizeof(netPluginLib_t));
|
||||
}
|
||||
return ncclSuccess;
|
||||
@@ -105,8 +105,9 @@ exit:
|
||||
return ncclSuccess;
|
||||
fail:
|
||||
if (pluginLib->dlHandle) {
|
||||
NCCLCHECK(ncclClosePluginLib(pluginLib->dlHandle));
|
||||
NCCLCHECK(ncclClosePluginLib(pluginLib->dlHandle, ncclPluginTypeNet));
|
||||
}
|
||||
pluginLib->dlHandle = nullptr;
|
||||
pluginLib->ncclNetPluginState = ncclNetPluginStateLoadFailed;
|
||||
pluginLib->ncclCollNetPluginState = ncclNetPluginStateLoadFailed;
|
||||
goto exit;
|
||||
|
||||
@@ -10,19 +10,15 @@
|
||||
#include <dlfcn.h>
|
||||
|
||||
#include "debug.h"
|
||||
#include "plugin.h"
|
||||
|
||||
#define MAX_STR_LEN 255
|
||||
|
||||
enum ncclPluginType {
|
||||
ncclPluginTypeNet,
|
||||
ncclPluginTypeTuner,
|
||||
ncclPluginTypeProfiler,
|
||||
};
|
||||
|
||||
#define NUM_LIBS 3
|
||||
static char* libNames[NUM_LIBS];
|
||||
static void *libHandles[NUM_LIBS];
|
||||
static const char *pluginNames[NUM_LIBS] = { "NET", "TUNER", "PROFILER" };
|
||||
static const char *pluginPrefix[NUM_LIBS] = { "librccl-net", "librccl-tuner", "librccl-profiler" };
|
||||
static const char *pluginPrefix[NUM_LIBS] = { "librccl-net", "libnccl-tuner", "libnccl-profiler" };
|
||||
static const char *pluginFallback[NUM_LIBS] = { "", "Using internal tuner plugin.", "" };
|
||||
static unsigned long subsys[NUM_LIBS] = { NCCL_INIT|NCCL_NET, NCCL_INIT|NCCL_TUNING, NCCL_INIT };
|
||||
|
||||
@@ -61,24 +57,26 @@ static void* openPluginLib(enum ncclPluginType type, const char* libName) {
|
||||
char eNoEntNameList[PATH_MAX] = { 0 };
|
||||
|
||||
if (libName && strlen(libName)) {
|
||||
// match names that start with 'lib' and end with '.so'
|
||||
if (strlen(libName) >= strlen("libX.so") && strncmp(libName, "lib", strlen("lib")) == 0 && strncmp(libName + strlen(libName) - strlen(".so"), ".so", strlen(".so")) == 0) {
|
||||
snprintf(libName_, MAX_STR_LEN, "%s", libName);
|
||||
libHandles[type] = tryOpenLib(libName_, &openErr, openErrStr);
|
||||
if (libHandles[type]) {
|
||||
INFO(subsys[type], "%s/Plugin: Plugin name set by env to %s", pluginNames[type], libName_);
|
||||
return libHandles[type];
|
||||
}
|
||||
if (openErr == ENOENT) {
|
||||
appendNameToList(eNoEntNameList, &len, libName_);
|
||||
} else {
|
||||
INFO(subsys[type], "%s/Plugin: %s", pluginNames[type], openErrStr);
|
||||
}
|
||||
snprintf(libName_, MAX_STR_LEN, "%s", libName);
|
||||
libHandles[type] = tryOpenLib(libName_, &openErr, openErrStr);
|
||||
if (libHandles[type]) {
|
||||
INFO(subsys[type], "%s/Plugin: Plugin name set by env to %s", pluginNames[type], libName_);
|
||||
libNames[type] = strdup(libName_);
|
||||
return libHandles[type];
|
||||
}
|
||||
if (openErr == ENOENT) {
|
||||
appendNameToList(eNoEntNameList, &len, libName_);
|
||||
} else {
|
||||
INFO(subsys[type], "%s/Plugin: %s", pluginNames[type], openErrStr);
|
||||
}
|
||||
|
||||
// libName can't be a relative or absolute path (start with '.' or contain any '/'). It can't be a library name either (start with 'lib' or end with '.so')
|
||||
if (strchr(libName, '/') == nullptr && (strncmp(libName, "lib", strlen("lib")) || strlen(libName) < strlen(".so") || strncmp(libName + strlen(libName) - strlen(".so"), ".so", strlen(".so")))) {
|
||||
snprintf(libName_, MAX_STR_LEN, "%s-%s.so", pluginPrefix[type], libName);
|
||||
libHandles[type] = tryOpenLib(libName_, &openErr, openErrStr);
|
||||
if (libHandles[type]) {
|
||||
INFO(subsys[type], "%s/Plugin: Plugin name set by env to %s", pluginNames[type], libName_);
|
||||
libNames[type] = strdup(libName_);
|
||||
return libHandles[type];
|
||||
}
|
||||
if (openErr == ENOENT) {
|
||||
@@ -91,6 +89,7 @@ static void* openPluginLib(enum ncclPluginType type, const char* libName) {
|
||||
snprintf(libName_, MAX_STR_LEN, "%s.so", pluginPrefix[type]);
|
||||
libHandles[type] = tryOpenLib(libName_, &openErr, openErrStr);
|
||||
if (libHandles[type]) {
|
||||
libNames[type] = strdup(libName_);
|
||||
return libHandles[type];
|
||||
}
|
||||
if (openErr == ENOENT) {
|
||||
@@ -120,22 +119,21 @@ void* ncclOpenProfilerPluginLib(const char* name) {
|
||||
return openPluginLib(ncclPluginTypeProfiler, name);
|
||||
}
|
||||
|
||||
void* ncclGetNetPluginLib(void) {
|
||||
return libHandles[ncclPluginTypeNet];
|
||||
void* ncclGetNetPluginLib(enum ncclPluginType type) {
|
||||
if (libNames[ncclPluginTypeNet]) {
|
||||
// increment the reference counter of the net library
|
||||
libNames[type] = strdup(libNames[ncclPluginTypeNet]);
|
||||
libHandles[type] = dlopen(libNames[ncclPluginTypeNet], RTLD_NOW | RTLD_LOCAL);
|
||||
}
|
||||
return libHandles[type];
|
||||
}
|
||||
|
||||
ncclResult_t ncclClosePluginLib(void* handle) {
|
||||
bool found = false;
|
||||
for (int l=0; l<NUM_LIBS; l++) {
|
||||
if (libHandles[l] == handle) {
|
||||
libHandles[l] = nullptr;
|
||||
if (!found) {
|
||||
if (handle) {
|
||||
dlclose(handle);
|
||||
}
|
||||
found = true;
|
||||
}
|
||||
}
|
||||
ncclResult_t ncclClosePluginLib(void* handle, enum ncclPluginType type) {
|
||||
if (handle && libHandles[type] == handle) {
|
||||
dlclose(handle);
|
||||
libHandles[type] = nullptr;
|
||||
free(libNames[type]);
|
||||
libNames[type] = nullptr;
|
||||
}
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
@@ -79,7 +79,8 @@ exit:
|
||||
pthread_mutex_unlock(&profilerLock);
|
||||
return ncclSuccess;
|
||||
fail:
|
||||
if (profilerPluginLib) NCCLCHECK(ncclClosePluginLib(profilerPluginLib));
|
||||
if (profilerPluginLib) NCCLCHECK(ncclClosePluginLib(profilerPluginLib, ncclPluginTypeProfiler));
|
||||
profilerPluginLib = nullptr;
|
||||
profilerPluginStatus = profilerPluginLoadFailed;
|
||||
goto exit;
|
||||
}
|
||||
@@ -88,7 +89,7 @@ static ncclResult_t ncclProfilerPluginUnload(void) {
|
||||
pthread_mutex_lock(&profilerLock);
|
||||
if (0 == (--profilerPluginRefCount)) {
|
||||
INFO(NCCL_ENV, "PROFILER/Plugin: Closing profiler plugin %s", ncclProfiler->name);
|
||||
NCCLCHECK(ncclClosePluginLib(profilerPluginLib));
|
||||
NCCLCHECK(ncclClosePluginLib(profilerPluginLib, ncclPluginTypeProfiler));
|
||||
profilerPluginLib = nullptr;
|
||||
ncclProfiler = nullptr;
|
||||
profilerPluginStatus = profilerPluginLoadReady;
|
||||
|
||||
@@ -52,7 +52,7 @@ ncclResult_t ncclTunerPluginLoad(struct ncclComm* comm) {
|
||||
|
||||
tunerPluginLib = ncclOpenTunerPluginLib(ncclGetEnv("NCCL_TUNER_PLUGIN"));
|
||||
if (nullptr == tunerPluginLib) {
|
||||
tunerPluginLib = ncclGetNetPluginLib();
|
||||
tunerPluginLib = ncclGetNetPluginLib(ncclPluginTypeTuner);
|
||||
if (nullptr == tunerPluginLib) {
|
||||
goto fail;
|
||||
}
|
||||
@@ -78,6 +78,7 @@ exit:
|
||||
pthread_mutex_unlock(&tunerPluginLock);
|
||||
return ncclSuccess;
|
||||
fail:
|
||||
if (tunerPluginLib) NCCLCHECK(ncclClosePluginLib(tunerPluginLib, ncclPluginTypeTuner));
|
||||
tunerPluginLib = nullptr;
|
||||
status = tunerPluginLoadFailed;
|
||||
goto exit;
|
||||
@@ -87,7 +88,7 @@ ncclResult_t ncclTunerPluginUnload(struct ncclComm* comm) {
|
||||
pthread_mutex_lock(&tunerPluginLock);
|
||||
if (comm->tunerPluginLoaded && 0 == (--tunerPluginRefCount)) {
|
||||
INFO(NCCL_TUNING, "TUNER/Plugin: Closing tuner: '%s'", tunerSymbol->name);
|
||||
NCCLCHECK(ncclClosePluginLib(tunerPluginLib));
|
||||
NCCLCHECK(ncclClosePluginLib(tunerPluginLib, ncclPluginTypeTuner));
|
||||
tunerPluginLib = nullptr;
|
||||
tunerSymbol = nullptr;
|
||||
comm->tuner = nullptr;
|
||||
|
||||
@@ -497,7 +497,9 @@ static int ibvSpeeds[] = {
|
||||
14000, /* FDR */
|
||||
25000, /* EDR */
|
||||
50000, /* HDR */
|
||||
100000 /* NDR */ };
|
||||
100000, /* NDR */
|
||||
200000 /* XDR */
|
||||
};
|
||||
|
||||
static int firstBitSet(int val, int max) {
|
||||
int i = 0;
|
||||
@@ -653,12 +655,15 @@ ncclResult_t ncclIbInit(ncclDebugLogger_t logFunction, ncclProfilerCallback_t pr
|
||||
enum ncclIbProvider ibProvider = IB_PROVIDER_NONE;
|
||||
char dataDirectDevicePath[PATH_MAX];
|
||||
int dataDirectSupported = 0;
|
||||
int skipNetDevForDataDirect = 0;
|
||||
if (wrap_mlx5dv_is_supported(devices[d])) {
|
||||
ibProvider = IB_PROVIDER_MLX5;
|
||||
snprintf(dataDirectDevicePath, PATH_MAX, "/sys");
|
||||
if((ncclMlx5dvDmaBufCapable(context)) && (wrap_mlx5dv_get_data_direct_sysfs_path(context, dataDirectDevicePath + 4, PATH_MAX - 4) == ncclSuccess)) {
|
||||
INFO(NCCL_NET, "Data Direct DMA Interface is detected for device:%s", devices[d]->name);
|
||||
if(ncclParamIbDataDirect()) dataDirectSupported = 1;
|
||||
INFO(NCCL_INIT|NCCL_NET, "NET/IB: Data Direct DMA Interface is detected for device:%s", devices[d]->name);
|
||||
// Now check whether Data Direct has been disabled by the user
|
||||
if(ncclParamIbDataDirect() == 1) { dataDirectSupported = 1; skipNetDevForDataDirect = 1; }
|
||||
if(ncclParamIbDataDirect() == 2) { dataDirectSupported = 1; skipNetDevForDataDirect = 0; }
|
||||
}
|
||||
}
|
||||
int nPorts = 0;
|
||||
@@ -674,7 +679,8 @@ ncclResult_t ncclIbInit(ncclDebugLogger_t logFunction, ncclProfilerCallback_t pr
|
||||
continue;
|
||||
}
|
||||
for (int port_num = 1; port_num <= devAttr.phys_port_cnt; port_num++) {
|
||||
for (int dataDirect = 0; dataDirect < 1 + dataDirectSupported; ++dataDirect) {
|
||||
// dataDirect = 0 exposes the devices normally, dataDirect = 1 exposes the devices through direct NIC
|
||||
for (int dataDirect = skipNetDevForDataDirect; dataDirect < 1 + dataDirectSupported; ++dataDirect) {
|
||||
struct ibv_port_attr portAttr;
|
||||
if (ncclSuccess != wrap_ibv_query_port(context, port_num, &portAttr)) {
|
||||
WARN("NET/IB : Unable to query port_num %d", port_num);
|
||||
@@ -695,15 +701,18 @@ ncclResult_t ncclIbInit(ncclDebugLogger_t logFunction, ncclProfilerCallback_t pr
|
||||
ncclIbDevs[ncclNIbDevs].portAttr = portAttr;
|
||||
ncclIbDevs[ncclNIbDevs].portNum = port_num;
|
||||
ncclIbDevs[ncclNIbDevs].link = portAttr.link_layer;
|
||||
ncclIbDevs[ncclNIbDevs].speed = ncclIbSpeed(portAttr.active_speed) * ncclIbWidth(portAttr.active_width);
|
||||
if (portAttr.active_speed_ex)
|
||||
// A non-zero active_speed_ex indicates XDR rate (0x100) or higher
|
||||
ncclIbDevs[ncclNIbDevs].speed = ncclIbSpeed(portAttr.active_speed_ex) * ncclIbWidth(portAttr.active_width);
|
||||
else
|
||||
ncclIbDevs[ncclNIbDevs].speed = ncclIbSpeed(portAttr.active_speed) * ncclIbWidth(portAttr.active_width);
|
||||
ncclIbDevs[ncclNIbDevs].context = context;
|
||||
ncclIbDevs[ncclNIbDevs].pdRefs = 0;
|
||||
ncclIbDevs[ncclNIbDevs].pd = NULL;
|
||||
if (!dataDirect) {
|
||||
strncpy(ncclIbDevs[ncclNIbDevs].devName, devices[d]->name, MAXNAMESIZE);
|
||||
NCCLCHECKGOTO(ncclIbGetPciPath(ncclIbDevs[ncclNIbDevs].devName, &ncclIbDevs[ncclNIbDevs].pciPath, &ncclIbDevs[ncclNIbDevs].realPort), ret, fail);
|
||||
}
|
||||
else {
|
||||
} else {
|
||||
snprintf(ncclIbDevs[ncclNIbDevs].devName, MAXNAMESIZE, "%s_dma", devices[d]->name);
|
||||
NCCLCHECK(ncclCalloc(&ncclIbDevs[ncclNIbDevs].pciPath, PATH_MAX));
|
||||
strncpy(ncclIbDevs[ncclNIbDevs].pciPath, dataDirectDevicePath, PATH_MAX);
|
||||
|
||||
Reference in New Issue
Block a user