Wenkai Du
@@ -284,7 +284,7 @@ ncclResult_t ncclTopoAddNet(struct ncclXmlNode* xmlNet, struct ncclTopoSystem* s
|
||||
if (xmlGetAttrInt(xmlNet, "port", &net->net.port) != ncclSuccess) net->net.port = 0;
|
||||
if (xmlGetAttrInt(xmlNet, "gdr", &net->net.gdrSupport) != ncclSuccess) net->net.gdrSupport = 0;
|
||||
if (xmlGetAttrInt(xmlNet, "maxconn", &net->net.maxChannels) != ncclSuccess) net->net.maxChannels = MAXCHANNELS;
|
||||
if (xmlGetAttrInt(xmlNet, "coll", &net->net.collSupport) != ncclSuccess) net->net.collSupport = 0;
|
||||
if (ncclCollNet && xmlGetAttrInt(xmlNet, "coll", &net->net.collSupport) != ncclSuccess) net->net.collSupport = 0;
|
||||
ncclDebugNoWarn = 0;
|
||||
|
||||
NCCLCHECK(ncclTopoConnectNodes(nic, net, LINK_NET, net->net.width));
|
||||
|
||||
@@ -8,7 +8,8 @@ HIPCC = $(HIP_PATH)/bin/hipcc
|
||||
EXE = topo_expl
|
||||
CXXFLAGS = -g -O3 -Iinclude -I../../src/include -I../../src/graph/ -DTOPO_EXPL -DENABLE_TRACE
|
||||
|
||||
files = $(EXE).cpp model.cpp utils.cpp ../../src/graph/topo.cc ../../src/graph/rings.cc ../../src/graph/paths.cc ../../src/graph/trees.cc ../../src/graph/search.cc ../../src/graph/connect.cc
|
||||
files = $(EXE).cpp model.cpp utils.cpp ../../src/graph/topo.cc ../../src/graph/rings.cc ../../src/graph/paths.cc ../../src/graph/trees.cc \
|
||||
../../src/graph/search.cc ../../src/graph/connect.cc ../../src/graph/tuning.cc ../../src/graph/xml.cc ../../src/misc/nvmlwrap_stub.cc
|
||||
|
||||
all: $(EXE)
|
||||
|
||||
|
||||
@@ -23,397 +23,114 @@ THE SOFTWARE.
|
||||
#ifndef MODEL_H_
|
||||
#define MODEL_H_
|
||||
|
||||
class CpuDevices {
|
||||
private:
|
||||
char *cpuName;
|
||||
int interCpuWidth;
|
||||
int cpuPciWidth;
|
||||
int p2pPciWidth;
|
||||
|
||||
public:
|
||||
CpuDevices(const char *cpuname, const int intercpuwidth, const int cpupciwidth, const int p2ppciwidth) :
|
||||
cpuName((char *)cpuname), interCpuWidth(intercpuwidth), cpuPciWidth(cpupciwidth), p2pPciWidth(p2ppciwidth) {}
|
||||
|
||||
CpuDevices() : cpuName(0), interCpuWidth(0), cpuPciWidth(0), p2pPciWidth(0) {}
|
||||
|
||||
ncclResult_t getCpuWidths(char* name, int* interCpu, int* cpuPci, int* p2pPci) {
|
||||
strcpy(name, cpuName);
|
||||
*interCpu = interCpuWidth;
|
||||
*cpuPci = cpuPciWidth;
|
||||
*p2pPci = p2pPciWidth;
|
||||
return ncclSuccess;
|
||||
}
|
||||
};
|
||||
|
||||
class GpuDevices {
|
||||
private:
|
||||
int nGpus;
|
||||
uint64_t *busIds;
|
||||
char **gpuPciPaths;
|
||||
int *gpuNumaIds;
|
||||
int *connMatrix;
|
||||
|
||||
public:
|
||||
GpuDevices(const int ngpus, const uint64_t *busids, const char **gpupcipaths, const int *gpunumaids, const int *connmatrix) :
|
||||
nGpus(ngpus), busIds((uint64_t *)busids), gpuPciPaths((char **)gpupcipaths), gpuNumaIds((int *)gpunumaids), connMatrix((int *)connmatrix) {}
|
||||
|
||||
GpuDevices () : nGpus(0), busIds(0), gpuPciPaths(0), gpuNumaIds(0), connMatrix(0) {}
|
||||
|
||||
int getnDevs() { return nGpus; }
|
||||
|
||||
uint64_t getBusId(int dev) { return busIds[dev]; }
|
||||
|
||||
ncclResult_t getPciPath(char* busId, char** path) {
|
||||
char tempBusId[] = "0000:00:00.0";
|
||||
*path = (char *)malloc(PATH_MAX);
|
||||
int i;
|
||||
for (i = 0; i < nGpus; i++) {
|
||||
NCCLCHECK(int64ToBusId(busIds[i], tempBusId));
|
||||
if (strcmp(busId, tempBusId) == 0)
|
||||
break;
|
||||
}
|
||||
if (i < nGpus)
|
||||
strcpy(*path, gpuPciPaths[i]);
|
||||
else {
|
||||
WARN("Could not find real path of %s", busId);
|
||||
return ncclSystemError;
|
||||
}
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
int p2pCanConnect(int device1, int device2) {
|
||||
// connection matrix are 8 GPUs
|
||||
int dist = connMatrix[device1*8+device2];
|
||||
if (dist == 255)
|
||||
return 0;
|
||||
//if (dist%15 == 0 && dist/15 != 1) {
|
||||
// return 0;
|
||||
//}
|
||||
return 1;
|
||||
};
|
||||
|
||||
hipError_t getLinkTypeAndHopCount(int device1, int device2, uint32_t* linktype, uint32_t* hopcount) {
|
||||
// connection matrix are 8 GPUs
|
||||
int dist = connMatrix[device1*8+device2];
|
||||
|
||||
if (dist%15 == 0) {
|
||||
*linktype = 4;
|
||||
*hopcount = dist/15;
|
||||
}
|
||||
else if (dist%20 == 0) {
|
||||
*linktype = 2;
|
||||
*hopcount = dist/20;
|
||||
}
|
||||
else if (dist%36 == 0) {
|
||||
*linktype = 1;
|
||||
*hopcount = dist/36;
|
||||
}
|
||||
return hipSuccess;
|
||||
}
|
||||
|
||||
virtual int getNumaId(char *path) {
|
||||
int n;
|
||||
// search for all GPUs
|
||||
for (n = 0; n < nGpus; n++)
|
||||
if (strcmp(path, gpuPciPaths[n]) == 0)
|
||||
break;
|
||||
if (n < nGpus)
|
||||
return gpuNumaIds[n];
|
||||
return -1;
|
||||
}
|
||||
};
|
||||
|
||||
class NetDevices {
|
||||
private:
|
||||
int nNetDevs;
|
||||
char **netPciPaths;
|
||||
uint64_t *netGuids; // IB ports on same card share the same GUID
|
||||
int *netNumaIds;
|
||||
|
||||
public:
|
||||
NetDevices(const int nnetdevs, const char **netpcipaths, const uint64_t *netguids, const int *netnumaids) :
|
||||
nNetDevs(nnetdevs), netPciPaths((char **)netpcipaths), netGuids((uint64_t *)netguids), netNumaIds((int *)netnumaids) {}
|
||||
|
||||
NetDevices() : nNetDevs(0), netPciPaths(0), netGuids(0), netNumaIds(0) {}
|
||||
|
||||
int getnDevs() { return nNetDevs; }
|
||||
|
||||
ncclResult_t getPciPath(int dev, char** path) {
|
||||
*path = (char *)malloc(PATH_MAX);
|
||||
if (dev < nNetDevs)
|
||||
strcpy(*path, netPciPaths[dev]);
|
||||
else {
|
||||
WARN("Could not find real path of %d", dev);
|
||||
return ncclSystemError;
|
||||
}
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
virtual int getNumaId(char *path) {
|
||||
int n;
|
||||
// search for all NICs
|
||||
for (n = 0; n < nNetDevs; n++)
|
||||
if (strcmp(path, netPciPaths[n]) == 0)
|
||||
break;
|
||||
if (n < nNetDevs)
|
||||
return netNumaIds[n];
|
||||
return -1;
|
||||
}
|
||||
|
||||
uint64_t getIbGuid(char* path) {
|
||||
int n;
|
||||
for (n = 0; n < nNetDevs; n++)
|
||||
if (strcmp(path, netPciPaths[n]) == 0)
|
||||
break;
|
||||
if (n < nNetDevs)
|
||||
return netGuids[n];
|
||||
WARN("Invalid IB path %s", path);
|
||||
return 0;
|
||||
}
|
||||
};
|
||||
#include <vector>
|
||||
#include "topo.h"
|
||||
#include "xml.h"
|
||||
#include "utils.h"
|
||||
|
||||
class NodeModel {
|
||||
private:
|
||||
CpuDevices cpus;
|
||||
GpuDevices gpus;
|
||||
NetDevices netdevs;
|
||||
|
||||
public:
|
||||
int nodeId;
|
||||
int currRank;
|
||||
int firstRank;
|
||||
std::vector<struct ncclTopoSystem*> systems;
|
||||
uint64_t hostHash; // auto-generated
|
||||
uint64_t pidHash; // auto-generated
|
||||
char description[256];
|
||||
int nodeId;
|
||||
int firstRank;
|
||||
int currRank;
|
||||
|
||||
int rankToCudaDev(int rank) { return rank - firstRank; }
|
||||
|
||||
int getnGpus() { return gpus.getnDevs(); }
|
||||
|
||||
int getnNetDevs() { return netdevs.getnDevs(); }
|
||||
|
||||
ncclResult_t getGpuPciPath(char* busId, char** path) {
|
||||
return gpus.getPciPath(busId, path);
|
||||
NodeModel(const char *xml_file) {
|
||||
char filename[PATH_MAX];
|
||||
ssize_t count = readlink("/proc/self/exe", filename, PATH_MAX);
|
||||
while (--count > 0) {
|
||||
if (filename[count] == '/') {
|
||||
filename[count+1] = 0;
|
||||
break;
|
||||
}
|
||||
};
|
||||
strcat(filename, "models/");
|
||||
strcat(filename, xml_file);
|
||||
struct ncclTopoSystem* system;
|
||||
ncclTopoGetSystem(filename, &system);
|
||||
systems.push_back(system);
|
||||
for (int i=0; i<getNumGpus()-1; i++) {
|
||||
ncclTopoGetSystem(filename, &system);
|
||||
systems.push_back(system);
|
||||
}
|
||||
hostHash = ((uint64_t)rand() << 32) | rand();
|
||||
pidHash = ((uint64_t)rand() << 32) | rand();
|
||||
}
|
||||
|
||||
ncclResult_t getNetPciPath(int dev, char** path) {
|
||||
struct ncclTopoSystem* getSystem(int rank) { return systems[rank-firstRank]; }
|
||||
|
||||
return netdevs.getPciPath(dev, path);
|
||||
int getNumGpus() {
|
||||
return systems[0]->nodes[GPU].count;
|
||||
}
|
||||
|
||||
uint64_t getGpuBusId(int dev) {
|
||||
return gpus.getBusId(dev);
|
||||
int rankToCudaDev(int rank) {
|
||||
for (int i=0; i<getNumGpus(); i++) {
|
||||
if (rank == systems[0]->nodes[GPU].nodes[i].gpu.rank)
|
||||
return systems[0]->nodes[GPU].nodes[i].gpu.dev;
|
||||
}
|
||||
return -1;
|
||||
}
|
||||
|
||||
int p2pCanConnect(int device1, int device2) { return gpus.p2pCanConnect(device1, device2); }
|
||||
|
||||
hipError_t getLinkTypeAndHopCount(int device1, int device2, uint32_t* linktype, uint32_t* hopcount) {
|
||||
return gpus.getLinkTypeAndHopCount(device1, device2, linktype, hopcount);
|
||||
int64_t getGpuBusId(int rank) {
|
||||
for (int i=0; i<getNumGpus(); i++) {
|
||||
if (rank == systems[0]->nodes[GPU].nodes[i].gpu.rank)
|
||||
return systems[0]->nodes[GPU].nodes[i].id;
|
||||
}
|
||||
return -1;
|
||||
}
|
||||
|
||||
uint64_t getIbGuid(char* path) {
|
||||
return netdevs.getIbGuid(path);
|
||||
int busIdToCudaDev(int64_t busId) {
|
||||
for (int i=0; i<getNumGpus(); i++)
|
||||
if (systems[0]->nodes[GPU].nodes[i].id == busId)
|
||||
return systems[0]->nodes[GPU].nodes[i].gpu.dev;
|
||||
return -1;
|
||||
}
|
||||
|
||||
void setRanks() {
|
||||
for (int r=0; r<getNumGpus(); r++)
|
||||
for (int i=0; i<getNumGpus(); i++)
|
||||
systems[r]->nodes[GPU].nodes[i].gpu.rank += firstRank;
|
||||
}
|
||||
|
||||
int p2pCanConnect(int device1, int device2) { return 1; }
|
||||
int shmCanConnect(int device1, int device2) { return 1; }
|
||||
int netCanConnect(int device1, int device2) { return 1; }
|
||||
|
||||
virtual int getNumaId(char *path) {
|
||||
int numa = gpus.getNumaId(path);
|
||||
if (numa != -1) return numa;
|
||||
numa = netdevs.getNumaId(path);
|
||||
if (numa != -1) return numa;
|
||||
WARN("Invalid path %s for getNumaId", path);
|
||||
return 0;
|
||||
}
|
||||
|
||||
virtual ncclResult_t getCpuWidths(char* name, int* interCpu, int* cpuPci, int* p2pPci) {
|
||||
return cpus.getCpuWidths(name, interCpu, cpuPci, p2pPci);
|
||||
}
|
||||
|
||||
NodeModel(CpuDevices cpu, GpuDevices gpu, NetDevices net, const char *desc) :
|
||||
cpus(cpu), gpus(gpu), netdevs(net) {
|
||||
strncpy(description, desc, 256);
|
||||
}
|
||||
|
||||
NodeModel() {}
|
||||
|
||||
~NodeModel() {}
|
||||
};
|
||||
|
||||
class NetworkModel {
|
||||
private:
|
||||
int nNodes;
|
||||
int nRanks;
|
||||
NodeModel nodes[NCCL_TOPO_MAX_NODES];
|
||||
std::vector<NodeModel*> nodes;
|
||||
|
||||
public:
|
||||
void AddNode(NodeModel node) {
|
||||
nodes[nNodes] = node;
|
||||
nodes[nNodes].nodeId = nNodes;
|
||||
nodes[nNodes].firstRank = nRanks;
|
||||
nodes[nNodes].hostHash = ((uint64_t)rand() << 32) | rand();
|
||||
nodes[nNodes].pidHash = ((uint64_t)rand() << 32) | rand();
|
||||
nNodes++;
|
||||
nRanks += node.getnGpus();
|
||||
void AddNode(NodeModel* node) {
|
||||
node->nodeId = nodes.size();
|
||||
node->firstRank = nRanks;
|
||||
node->setRanks();
|
||||
nRanks += node->getNumGpus();
|
||||
nodes.push_back(node);
|
||||
}
|
||||
|
||||
int GetNNodes() { return nNodes; }
|
||||
|
||||
int GetNRanks() { return nRanks; }
|
||||
|
||||
NodeModel* GetNode(int rank) {
|
||||
int node_id;
|
||||
|
||||
if(rank < 0 || rank >= nRanks)
|
||||
return 0;
|
||||
|
||||
for(node_id = nNodes-1; node_id >= 0; node_id--)
|
||||
if(rank >= nodes[node_id].firstRank) break;
|
||||
|
||||
if (node_id >= 0) {
|
||||
nodes[node_id].currRank = rank;
|
||||
return nodes+node_id;
|
||||
for (auto & node : nodes) {
|
||||
if (rank >= node->firstRank && rank < node->firstRank+node->getNumGpus()) {
|
||||
node->currRank = rank;
|
||||
return node;
|
||||
}
|
||||
}
|
||||
else
|
||||
return 0;
|
||||
return NULL;
|
||||
}
|
||||
|
||||
NetworkModel() : nNodes(0), nRanks(0) {}
|
||||
};
|
||||
int GetNNodes() { return nodes.size(); }
|
||||
int GetNRanks() { return nRanks; }
|
||||
|
||||
|
||||
const static uint64_t busIds_8[] = { 0x1d000, 0x20000, 0x23000, 0x26000, 0x3f000, 0x43000, 0x46000, 0x49000 };
|
||||
|
||||
const static char* gpuPciPaths_8[] = {
|
||||
"/sys/devices/pci0000:17/0000:17:00.0/0000:18:00.0/0000:19:08.0/0000:1b:00.0/0000:1c:00.0/0000:1d:00.0",
|
||||
"/sys/devices/pci0000:17/0000:17:00.0/0000:18:00.0/0000:19:0c.0/0000:1e:00.0/0000:1f:00.0/0000:20:00.0",
|
||||
"/sys/devices/pci0000:17/0000:17:00.0/0000:18:00.0/0000:19:10.0/0000:21:00.0/0000:22:00.0/0000:23:00.0",
|
||||
"/sys/devices/pci0000:17/0000:17:00.0/0000:18:00.0/0000:19:14.0/0000:24:00.0/0000:25:00.0/0000:26:00.0",
|
||||
"/sys/devices/pci0000:3a/0000:3a:00.0/0000:3b:00.0/0000:3c:04.0/0000:3d:00.0/0000:3e:00.0/0000:3f:00.0",
|
||||
"/sys/devices/pci0000:3a/0000:3a:00.0/0000:3b:00.0/0000:3c:0c.0/0000:41:00.0/0000:42:00.0/0000:43:00.0",
|
||||
"/sys/devices/pci0000:3a/0000:3a:00.0/0000:3b:00.0/0000:3c:10.0/0000:44:00.0/0000:45:00.0/0000:46:00.0",
|
||||
"/sys/devices/pci0000:3a/0000:3a:00.0/0000:3b:00.0/0000:3c:14.0/0000:47:00.0/0000:48:00.0/0000:49:00.0",
|
||||
};
|
||||
|
||||
const static int gpuPciNumaIds_8[] = { 0, 0, 0, 0, 0, 0, 0, 0 };
|
||||
|
||||
const static char* netPciPaths_1[] = {
|
||||
"/sys/devices/pci0000:17/0000:17:00.0/0000:18:00.0/0000:19:04.0/0000:1a:00.0",
|
||||
};
|
||||
|
||||
const static char* netPciPaths_1_1[] = {
|
||||
"/sys/devices/pci0000:3a/0000:3a:00.0/0000:3b:00.0/0000:3c:08.0/0000:4c:00.0",
|
||||
};
|
||||
|
||||
const static uint64_t netGuids_1[] = {
|
||||
0xb8599f030007053aL,
|
||||
};
|
||||
|
||||
const static int netPciNumaIds_1[] = { 0 };
|
||||
|
||||
const static char* netPciPaths_2[] = {
|
||||
"/sys/devices/pci0000:17/0000:17:00.0/0000:18:00.0/0000:19:04.0/0000:1a:00.0",
|
||||
"/sys/devices/pci0000:3a/0000:3a:00.0/0000:3b:00.0/0000:3c:08.0/0000:4c:00.0",
|
||||
};
|
||||
|
||||
const static uint64_t netGuids_2[] = {
|
||||
0xb8599f030007053aL,
|
||||
0x506b4b030027bbf2L,
|
||||
};
|
||||
|
||||
const static int netPciNumaIds_2[] = { 0, 0 };
|
||||
|
||||
const static uint64_t rome_busIds_8[] = { 0x63000, 0x23000, 0x26000, 0x03000, 0xe3000, 0xc3000, 0xc6000, 0xa3000 };
|
||||
|
||||
const static char* rome_gpuPciPaths_8[] = {
|
||||
"/sys/devices/pci0000:60/0000:60:03.1/0000:61:00.0/0000:62:00.0/0000:63:00.0",
|
||||
"/sys/devices/pci0000:20/0000:20:01.1/0000:21:00.0/0000:22:00.0/0000:23:00.0",
|
||||
"/sys/devices/pci0000:20/0000:20:03.1/0000:24:00.0/0000:25:00.0/0000:26:00.0",
|
||||
"/sys/devices/pci0000:00/0000:00:01.1/0000:01:00.0/0000:02:00.0/0000:03:00.0",
|
||||
"/sys/devices/pci0000:e0/0000:e0:03.1/0000:e1:00.0/0000:e2:00.0/0000:e3:00.0",
|
||||
"/sys/devices/pci0000:c0/0000:c0:01.1/0000:c1:00.0/0000:c2:00.0/0000:c3:00.0",
|
||||
"/sys/devices/pci0000:c0/0000:c0:03.1/0000:c4:00.0/0000:c5:00.0/0000:c6:00.0",
|
||||
"/sys/devices/pci0000:a0/0000:a0:03.1/0000:a1:00.0/0000:a2:00.0/0000:a3:00.0",
|
||||
};
|
||||
|
||||
const static int rome_gpuPciNumaIds_8[] = { 0, 0, 0, 0, 4, 4, 4, 4 };
|
||||
|
||||
const static char* rome_netPciPaths_1[] = {
|
||||
"/sys/devices/pci0000:40/0000:40:01.1/0000:41:00.0",
|
||||
};
|
||||
|
||||
const static uint64_t rome_netGuids_1[] = {
|
||||
0xb8599f030007053aL,
|
||||
};
|
||||
|
||||
const static int rom_netPciNumaIds_1[] = { 0 };
|
||||
|
||||
const static char* rome_netPciPaths_2[] = {
|
||||
"/sys/devices/pci0000:40/0000:40:01.1/0000:41:00.0",
|
||||
"/sys/devices/pci0000:80/0000:80:01.1/0000:81:00.0",
|
||||
};
|
||||
|
||||
const static uint64_t rome_netGuids_2[] = {
|
||||
0xb8599f030007053aL,
|
||||
0x506b4b030027bbf2L,
|
||||
};
|
||||
|
||||
const static int rom_netPciNumaIds_2[] = { 0, 4 };
|
||||
|
||||
const int conn_mat_pcie[64] = {
|
||||
0 , 40, 40, 40, 40, 40, 40, 40,
|
||||
40, 0 , 40, 40, 40, 40, 40, 40,
|
||||
40, 40, 0 , 40, 40, 40, 40, 40,
|
||||
40, 40, 40, 0 , 40, 40, 40, 40,
|
||||
40, 40, 40, 40, 0 , 40, 40, 40,
|
||||
40, 40, 40, 40, 40, 0 , 40, 40,
|
||||
40, 40, 40, 40, 40, 40, 0 , 40,
|
||||
40, 40, 40, 40, 40, 40, 40, 0 ,
|
||||
};
|
||||
|
||||
const int conn_mat_4p2h[64] = {
|
||||
0 , 15, 15, 30, 40, 40, 40, 40,
|
||||
15, 0 , 30, 15, 40, 40, 40, 40,
|
||||
15, 30, 0 , 15, 40, 40, 40, 40,
|
||||
30, 15, 15, 0 , 40, 40, 40, 40,
|
||||
40, 40, 40, 40, 0 , 15, 15, 30,
|
||||
40, 40, 40, 40, 15, 0 , 30, 15,
|
||||
40, 40, 40, 40, 15, 30, 0 , 15,
|
||||
40, 40, 40, 40, 30, 15, 15, 0 ,
|
||||
};
|
||||
|
||||
const int conn_mat_8p6l[64] = {
|
||||
0 , 15, 15, 15, 15, 30, 15, 15,
|
||||
15, 0 , 15, 15, 30, 15, 15, 15,
|
||||
15, 15, 0 , 15, 15, 15, 15, 30,
|
||||
15, 15, 15, 0 , 15, 15, 30, 15,
|
||||
15, 30, 15, 15, 0 , 15, 15, 15,
|
||||
30, 15, 15, 15, 15, 0 , 15, 15,
|
||||
15, 15, 15, 30, 15, 15, 0 , 15,
|
||||
15, 15, 30, 15, 15, 15, 15, 0 ,
|
||||
};
|
||||
|
||||
const int conn_mat_8p6l_1[64] = {
|
||||
0 , 15, 15, 30, 15, 15, 15, 15,
|
||||
15, 0 , 30, 15, 15, 15, 15, 15,
|
||||
15, 30, 0 , 15, 15, 15, 15, 15,
|
||||
30, 15, 15, 0 , 15, 15, 15, 15,
|
||||
15, 15, 15, 15, 0 , 15, 15, 30,
|
||||
15, 15, 15, 15, 15, 0 , 30, 15,
|
||||
15, 15, 15, 15, 15, 30, 0 , 15,
|
||||
15, 15, 15, 15, 30, 15, 15, 0 ,
|
||||
};
|
||||
|
||||
const int conn_mat_rome[64] = {
|
||||
0 , 40, 40, 40, 72, 72, 72, 72,
|
||||
40, 0 , 40, 40, 72, 72, 72, 72,
|
||||
40, 40, 0 , 40, 72, 72, 72, 72,
|
||||
40, 40, 40, 0 , 72, 72, 72, 72,
|
||||
72, 72, 72, 72, 0 , 40, 40, 40,
|
||||
72, 72, 72, 72, 40, 0 , 40, 40,
|
||||
72, 72, 72, 72, 40, 40, 0 , 40,
|
||||
72, 72, 72, 72, 40, 40, 40, 0 ,
|
||||
NetworkModel() : nRanks(0) {}
|
||||
};
|
||||
|
||||
#endif
|
||||
@@ -12,11 +12,11 @@
|
||||
#include <hip/hip_fp16.h>
|
||||
|
||||
#define NCCL_MAJOR 2
|
||||
#define NCCL_MINOR 5
|
||||
#define NCCL_PATCH 6
|
||||
#define NCCL_MINOR 6
|
||||
#define NCCL_PATCH 2
|
||||
#define NCCL_SUFFIX ""
|
||||
|
||||
#define NCCL_VERSION_CODE 2506
|
||||
#define NCCL_VERSION_CODE 2602
|
||||
#define NCCL_VERSION(X,Y,Z) ((X) * 1000 + (Y) * 100 + (Z))
|
||||
|
||||
#define RCCL_BFLOAT16 1
|
||||
|
||||
@@ -13,32 +13,34 @@ struct allGather1Data_t {
|
||||
struct ncclComm* comm;
|
||||
};
|
||||
|
||||
struct allGather3Data_t {
|
||||
// AllGather3 - begin
|
||||
struct ncclGraphInfo {
|
||||
int sameChannels;
|
||||
float speedIntra;
|
||||
float speedInter;
|
||||
int typeIntra;
|
||||
};
|
||||
|
||||
struct allGather3Data_t{
|
||||
int cudaCompCap;
|
||||
int fullCudaCompCap;
|
||||
int nvlink;
|
||||
int nChannels;
|
||||
struct {
|
||||
int sameChannels;
|
||||
int speedIntra;
|
||||
int speedInter;
|
||||
int nvlink;
|
||||
} tree;
|
||||
struct {
|
||||
int sameChannels;
|
||||
int speedIntra;
|
||||
int speedInter;
|
||||
int nvlink;
|
||||
} ring;
|
||||
struct ncclGraphInfo tree;
|
||||
struct ncclGraphInfo ring;
|
||||
struct ncclGraphInfo collNet;
|
||||
struct ncclTopoRanks topoRanks;
|
||||
};
|
||||
|
||||
ncclResult_t ncclTopoGetSystem(const char* xmlTopoFile, struct ncclTopoSystem** system);
|
||||
|
||||
ncclResult_t ncclTopoGetSystemFromXml(struct ncclXml* xml, struct ncclTopoSystem** topoSystem);
|
||||
|
||||
ncclResult_t bootstrapAllGather(struct ncclComm* comm, struct allGather1Data_t * allGather1Data);
|
||||
|
||||
ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather1Data_t *allGather1Data,
|
||||
struct allGather3Data_t *allGather3Data, struct ncclTopoGraph& treeGraph, struct ncclTopoGraph& ringGraph);
|
||||
ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather1Data_t *allGather1Data, struct allGather3Data_t *allGather3Data,
|
||||
struct ncclTopoGraph& treeGraph, struct ncclTopoGraph& ringGraph, struct ncclTopoGraph& collNetGraph);
|
||||
|
||||
ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t *allGather3Data,
|
||||
struct ncclTopoGraph& treeGraph, struct ncclTopoGraph& ringGraph);
|
||||
struct ncclTopoGraph& treeGraph, struct ncclTopoGraph& ringGraph, struct ncclTopoGraph& collNetGraph);
|
||||
|
||||
#endif
|
||||
@@ -41,24 +41,14 @@ THE SOFTWARE.
|
||||
#include <sys/stat.h>
|
||||
#include <unistd.h>
|
||||
#include "model.h"
|
||||
#include "topo.h"
|
||||
|
||||
extern NodeModel *node_model;
|
||||
|
||||
static ncclResult_t dummyNetDevices(int* ndev) {
|
||||
*ndev = node_model->getnNetDevs();
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
static ncclResult_t dummyNetPciPath(int dev, char** path) {
|
||||
node_model->getNetPciPath(dev, path);
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
ncclNet_t ncclNetDummy = {
|
||||
"IB",
|
||||
0,
|
||||
dummyNetDevices,
|
||||
dummyNetPciPath,
|
||||
0,
|
||||
0,
|
||||
0,
|
||||
0,
|
||||
@@ -76,24 +66,9 @@ ncclNet_t ncclNetDummy = {
|
||||
|
||||
ncclNet_t* ncclNet = &ncclNetDummy;
|
||||
|
||||
ncclResult_t wrapNvmlDeviceGetHandleByPciBusId(const char* pciBusId, nvmlDevice_t* device) {
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
/* Convert a PCI busId string into a local cudaDev device index (cf. CUDA_VISIBLE_DEVICES) */
|
||||
int busIdToCudaDev(int64_t busId) {
|
||||
int cudaDev;
|
||||
|
||||
for (cudaDev = 0; cudaDev < node_model->getnGpus(); cudaDev++) {
|
||||
if (node_model->getGpuBusId(cudaDev) == busId)
|
||||
break;
|
||||
}
|
||||
|
||||
if (cudaDev < node_model->getnGpus())
|
||||
return cudaDev;
|
||||
else
|
||||
WARN("Invalid busId %lx", busId);
|
||||
return 0;
|
||||
return node_model->busIdToCudaDev(busId);
|
||||
}
|
||||
|
||||
/* Determine if two peers can communicate with P2P */
|
||||
@@ -177,6 +152,8 @@ ncclResult_t netSendSetup(struct ncclTopoSystem* topo, struct ncclTopoGraph* gra
|
||||
int netDev, useGdr = 0;
|
||||
|
||||
NCCLCHECK(ncclTopoGetNetDev(graph, 1, channelId, &netDev));
|
||||
NCCLCHECK(ncclTopoCheckGdr(topo, myInfo->busId, netDev, 1, &useGdr));
|
||||
|
||||
INFO(NCCL_INIT|NCCL_NET,"Ring %02d : %d[%lx] -> %d[%lx] [send] via NET/%s/%d%s", channelId, myInfo->rank, myInfo->busId, peerInfo->rank, peerInfo->busId, ncclNetName(), netDev,
|
||||
useGdr ? "/GDRDMA" : "");
|
||||
return ncclSuccess;
|
||||
@@ -188,15 +165,8 @@ ncclResult_t netRecvSetup(struct ncclTopoSystem* topo, struct ncclTopoGraph* gra
|
||||
int netDev, useGdr = 0;
|
||||
|
||||
NCCLCHECK(ncclTopoGetNetDev(graph, 0, channelId, &netDev));
|
||||
// Check if we are close enough that it makes sense to enable GDR
|
||||
int netGdrLevel = ncclParamNetGdrLevel();
|
||||
int distance;
|
||||
NCCLCHECK(ncclTopoNetDistance(topo, myInfo->busId, netDev, &distance));
|
||||
if (distance >= netGdrLevel) {
|
||||
INFO(NCCL_NET,"NET/%s : GPU Direct RDMA Disabled for GPU %lx / HCA %d (distance %d >= %d)", ncclNetName(), myInfo->busId, netDev, distance, netGdrLevel);
|
||||
}
|
||||
else
|
||||
useGdr = 1;
|
||||
NCCLCHECK(ncclTopoCheckGdr(topo, myInfo->busId, netDev, 0, &useGdr));
|
||||
|
||||
INFO(NCCL_INIT|NCCL_NET,"Ring %02d : %d[%lx] -> %d[%lx] [receive] via NET/%s/%d%s", channelId, peerInfo->rank, peerInfo->busId, myInfo->rank, myInfo->busId, ncclNetName(), netDev,
|
||||
useGdr ? "/GDRDMA" : "");
|
||||
return ncclSuccess;
|
||||
|
||||
@@ -0,0 +1,43 @@
|
||||
<system version="1">
|
||||
<cpu numaid="0" affinity="0003ff,f0003fff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
|
||||
<pci busid="0000:18:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:1b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:1d:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="0" sm="30" rank="0" gdr="1">
|
||||
<xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:1e:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:20:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="1" sm="30" rank="1" gdr="1">
|
||||
<xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:21:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:23:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="2" sm="30" rank="2" gdr="1">
|
||||
<xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:24:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:26:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="3" sm="30" rank="3" gdr="1">
|
||||
<xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:1a:00.0" class="0x020000" link_speed="8 GT/s" link_width="16">
|
||||
<nic>
|
||||
<net name="mlx5_0" dev="0" speed="100000" port="1" guid="0xf2bb2700034b6b50" maxconn="262144" gdr="1"/>
|
||||
</nic>
|
||||
</pci>
|
||||
</pci>
|
||||
</cpu>
|
||||
</system>
|
||||
@@ -0,0 +1,46 @@
|
||||
<system version="1">
|
||||
<cpu numaid="0" affinity="0003ff,f0003fff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
|
||||
<pci busid="0000:18:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:1a:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="0" sm="30" rank="0" gdr="0">
|
||||
<xgmi target="0000:3d:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:b1:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:3b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:3d:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="1" sm="30" rank="1" gdr="0">
|
||||
<xgmi target="0000:1a:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:88:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:01:00.0" class="0x020000" link_speed="8 GT/s" link_width="4">
|
||||
<nic>
|
||||
<net name="eno1" dev="0" speed="10000" port="0" guid="0x0" maxconn="65536" gdr="0"/>
|
||||
</nic>
|
||||
</pci>
|
||||
<nic>
|
||||
<net name="virbr0" dev="1" speed="10000" port="0" guid="0x1" maxconn="65536" gdr="0"/>
|
||||
</nic>
|
||||
</cpu>
|
||||
<cpu numaid="1" affinity="fffc00,0fffc000" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
|
||||
<pci busid="0000:86:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:88:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="2" sm="30" rank="2" gdr="0">
|
||||
<xgmi target="0000:3d:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:b1:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:af:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:b1:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="3" sm="30" rank="3" gdr="0">
|
||||
<xgmi target="0000:1a:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:88:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
</cpu>
|
||||
</system>
|
||||
@@ -0,0 +1,77 @@
|
||||
<system version="1">
|
||||
<cpu numaid="0" affinity="0003ff,f0003fff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
|
||||
<pci busid="0000:18:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:1b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:1d:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="0" sm="30" rank="0" gdr="1">
|
||||
<xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:1e:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:20:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="1" sm="30" rank="1" gdr="1">
|
||||
<xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:21:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:23:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="2" sm="30" rank="2" gdr="1">
|
||||
<xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:24:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:26:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="3" sm="30" rank="3" gdr="1">
|
||||
<xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:1a:00.0" class="0x020000" link_speed="8 GT/s" link_width="16">
|
||||
<nic>
|
||||
<net name="mlx5_0" dev="0" speed="100000" port="1" guid="0xf2bb2700034b6b50" maxconn="262144" gdr="1"/>
|
||||
</nic>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:3b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:3d:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:3f:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="4" sm="30" rank="4" gdr="1">
|
||||
<xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:46:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:41:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:43:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="5" sm="30" rank="5" gdr="1">
|
||||
<xgmi target="0000:3f:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:49:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:44:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:46:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="6" sm="30" rank="6" gdr="1">
|
||||
<xgmi target="0000:3f:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:49:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:47:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:49:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="7" sm="30" rank="7" gdr="1">
|
||||
<xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:46:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
</cpu>
|
||||
</system>
|
||||
@@ -0,0 +1,77 @@
|
||||
<system version="1">
|
||||
<cpu numaid="0" affinity="0003ff,f0003fff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
|
||||
<pci busid="0000:18:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:1b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:1d:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="0" sm="30" rank="0" gdr="1">
|
||||
<xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:1e:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:20:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="1" sm="30" rank="1" gdr="1">
|
||||
<xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:21:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:23:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="2" sm="30" rank="2" gdr="1">
|
||||
<xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:24:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:26:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="3" sm="30" rank="3" gdr="1">
|
||||
<xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:3b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:3d:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:3f:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="4" sm="30" rank="4" gdr="1">
|
||||
<xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:46:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:41:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:43:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="5" sm="30" rank="5" gdr="1">
|
||||
<xgmi target="0000:3f:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:49:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:44:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:46:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="6" sm="30" rank="6" gdr="1">
|
||||
<xgmi target="0000:3f:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:49:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:47:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:49:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="7" sm="30" rank="7" gdr="1">
|
||||
<xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:46:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:4c:00.0" class="0x020000" link_speed="8 GT/s" link_width="16">
|
||||
<nic>
|
||||
<net name="mlx5_0" dev="0" speed="100000" port="1" guid="0xf2bb2700034b6b50" maxconn="262144" gdr="1"/>
|
||||
</nic>
|
||||
</pci>
|
||||
</pci>
|
||||
</cpu>
|
||||
</system>
|
||||
@@ -0,0 +1,82 @@
|
||||
<system version="1">
|
||||
<cpu numaid="0" affinity="0003ff,f0003fff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
|
||||
<pci busid="0000:18:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:1b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:1d:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="0" sm="30" rank="0" gdr="1">
|
||||
<xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:1e:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:20:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="1" sm="30" rank="1" gdr="1">
|
||||
<xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:21:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:23:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="2" sm="30" rank="2" gdr="1">
|
||||
<xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:24:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:26:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="3" sm="30" rank="3" gdr="1">
|
||||
<xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:1a:00.0" class="0x020000" link_speed="8 GT/s" link_width="16">
|
||||
<nic>
|
||||
<net name="mlx5_0" dev="0" speed="100000" port="1" guid="0xf2bb2700034b6b50" maxconn="262144" gdr="1"/>
|
||||
</nic>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:3b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:3d:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:3f:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="4" sm="30" rank="4" gdr="1">
|
||||
<xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:46:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:41:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:43:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="5" sm="30" rank="5" gdr="1">
|
||||
<xgmi target="0000:3f:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:49:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:44:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:46:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="6" sm="30" rank="6" gdr="1">
|
||||
<xgmi target="0000:3f:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:49:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:47:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:49:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="7" sm="30" rank="7" gdr="1">
|
||||
<xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:46:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:4c:00.0" class="0x020700" link_speed="8 GT/s" link_width="16">
|
||||
<nic>
|
||||
<net name="mlx5_1" dev="1" speed="100000" port="1" guid="0xb8599f030007053a" maxconn="262144" gdr="1"/>
|
||||
</nic>
|
||||
</pci>
|
||||
</pci>
|
||||
</cpu>
|
||||
</system>
|
||||
@@ -0,0 +1,52 @@
|
||||
<system version="1">
|
||||
<cpu numaid="0" affinity="0000ffff,0000ffff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
|
||||
<pci busid="0000:18:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:1b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:1d:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="0" sm="30" rank="0" gdr="0">
|
||||
<xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:1e:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:20:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="1" sm="30" rank="1" gdr="0">
|
||||
<xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:21:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:23:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="2" sm="30" rank="2" gdr="0">
|
||||
<xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:26:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:24:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:26:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="3" sm="30" rank="3" gdr="0">
|
||||
<xgmi target="0000:1d:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:20:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:23:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:5e:00.0" class="0x060400" link_speed="8 GT/s" link_width="8">
|
||||
<pci busid="0000:60:00.0" class="0x020000" link_speed="2.5 GT/s" link_width="1">
|
||||
<nic>
|
||||
<net name="enp96s0f0" dev="0" speed="10000" port="0" guid="0x0" maxconn="65536" gdr="0"/>
|
||||
</nic>
|
||||
</pci>
|
||||
</pci>
|
||||
<nic>
|
||||
<net name="veth608058c" dev="1" speed="10000" port="0" guid="0x1" maxconn="65536" gdr="0"/>
|
||||
</nic>
|
||||
</cpu>
|
||||
</system>
|
||||
@@ -0,0 +1,107 @@
|
||||
<system version="1">
|
||||
<cpu numaid="0" affinity="00000000,00000000,00000000,00000000,00000000,000000ff,ffff0000,00ffffff" arch="x86_64" vendor="AuthenticAMD" familyid="143" modelid="49">
|
||||
<pci busid="0000:61:00.0" class="0x060400" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:63:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="0" sm="30" rank="0" gdr="0">
|
||||
<xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:24:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:03:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:e3:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:a3:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:83:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:41:00.0" class="0x060400" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:43:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="1" sm="30" rank="1" gdr="0">
|
||||
<xgmi target="0000:63:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:24:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:03:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:c3:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:a3:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:83:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:22:00.0" class="0x060400" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:24:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="2" sm="30" rank="2" gdr="0">
|
||||
<xgmi target="0000:63:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:03:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:e3:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:c3:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:a3:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:01:00.0" class="0x060400" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:03:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="3" sm="30" rank="3" gdr="0">
|
||||
<xgmi target="0000:63:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:24:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:e3:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:c3:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:83:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
</cpu>
|
||||
<cpu numaid="1" affinity="00000000,00000000,00000000,00000000,00000000,ffffff00,0000ffff,ff000000" arch="x86_64" vendor="AuthenticAMD" familyid="143" modelid="49">
|
||||
<pci busid="0000:e1:00.0" class="0x060400" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:e3:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="4" sm="30" rank="4" gdr="0">
|
||||
<xgmi target="0000:63:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:24:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:03:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:c3:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:a3:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:83:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:c1:00.0" class="0x060400" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:c3:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="5" sm="30" rank="5" gdr="0">
|
||||
<xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:24:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:03:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:e3:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:a3:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:83:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:a1:00.0" class="0x060400" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:a3:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="6" sm="30" rank="6" gdr="0">
|
||||
<xgmi target="0000:63:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:24:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:e3:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:c3:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:83:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:81:00.0" class="0x060400" link_speed="16 GT/s" link_width="16">
|
||||
<pci busid="0000:83:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="7" sm="30" rank="7" gdr="0">
|
||||
<xgmi target="0000:63:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:43:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:03:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:e3:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:c3:00.0" count="1" tclass="0x038000"/>
|
||||
<xgmi target="0000:a3:00.0" count="1" tclass="0x038000"/>
|
||||
</gpu>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:c4:00.0" class="0x020000" link_speed="2.5 GT/s" link_width="1">
|
||||
<nic>
|
||||
<net name="enp196s0" dev="0" speed="10000" port="0" guid="0x0" maxconn="65536" gdr="0"/>
|
||||
</nic>
|
||||
</pci>
|
||||
</cpu>
|
||||
</system>
|
||||
@@ -0,0 +1,53 @@
|
||||
<system version="1">
|
||||
<cpu numaid="0" affinity="00ff00ff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
|
||||
<pci busid="0000:18:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:1b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:1d:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="0" sm="30" rank="0" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:1e:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:20:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="1" sm="30" rank="1" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:21:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:23:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="2" sm="30" rank="2" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:24:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:26:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="3" sm="30" rank="3" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:1a:00.0" class="0x020700" link_speed="8 GT/s" link_width="16">
|
||||
<nic>
|
||||
<net name="mlx5_0" dev="0" speed="100000" port="1" guid="0x3a050700039f59b8" maxconn="262144" gdr="1"/>
|
||||
</nic>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:3b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:3d:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:3f:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="4" sm="30" rank="4" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:41:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:43:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="5" sm="30" rank="5" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:44:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:46:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="6" sm="30" rank="6" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:47:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:49:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="7" sm="30" rank="7" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
</cpu>
|
||||
</system>
|
||||
@@ -0,0 +1,53 @@
|
||||
<system version="1">
|
||||
<cpu numaid="0" affinity="00ff00ff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
|
||||
<pci busid="0000:18:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:1b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:1d:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="0" sm="30" rank="0" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:1e:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:20:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="1" sm="30" rank="1" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:21:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:23:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="2" sm="30" rank="2" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:24:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:26:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="3" sm="30" rank="3" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:3b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:3d:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:3f:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="4" sm="30" rank="4" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:41:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:43:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="5" sm="30" rank="5" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:44:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:46:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="6" sm="30" rank="6" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:47:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:49:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="7" sm="30" rank="7" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:4c:00.0" class="0x020700" link_speed="8 GT/s" link_width="16">
|
||||
<nic>
|
||||
<net name="mlx5_0" dev="0" speed="100000" port="1" guid="0xb8599f030007053a" maxconn="262144" gdr="1"/>
|
||||
</nic>
|
||||
</pci>
|
||||
</pci>
|
||||
</cpu>
|
||||
</system>
|
||||
@@ -0,0 +1,58 @@
|
||||
<system version="1">
|
||||
<cpu numaid="0" affinity="00ff00ff" arch="x86_64" vendor="GenuineIntel" familyid="6" modelid="85">
|
||||
<pci busid="0000:18:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:1b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:1d:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="0" sm="30" rank="0" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:1e:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:20:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="1" sm="30" rank="1" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:21:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:23:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="2" sm="30" rank="2" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:24:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:26:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="3" sm="30" rank="3" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:1a:00.0" class="0x020700" link_speed="8 GT/s" link_width="16">
|
||||
<nic>
|
||||
<net name="mlx5_0" dev="0" speed="100000" port="1" guid="0x3a050700039f59b8" maxconn="262144" gdr="1"/>
|
||||
</nic>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:3b:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:3d:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:3f:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="4" sm="30" rank="4" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:41:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:43:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="5" sm="30" rank="5" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:44:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:46:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="6" sm="30" rank="6" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:47:00.0" class="0x060400" link_speed="8 GT/s" link_width="16">
|
||||
<pci busid="0000:49:00.0" class="0x038000" link_speed="16 GT/s" link_width="16">
|
||||
<gpu dev="7" sm="30" rank="7" gdr="1"/>
|
||||
</pci>
|
||||
</pci>
|
||||
<pci busid="0000:4c:00.0" class="0x020700" link_speed="8 GT/s" link_width="16">
|
||||
<nic>
|
||||
<net name="mlx5_1" dev="1" speed="100000" port="1" guid="0xb8599f030007053a" maxconn="262144" gdr="1"/>
|
||||
</nic>
|
||||
</pci>
|
||||
</pci>
|
||||
</cpu>
|
||||
</system>
|
||||
@@ -29,7 +29,6 @@ THE SOFTWARE.
|
||||
#include "net.h"
|
||||
#include "graph.h"
|
||||
#include "argcheck.h"
|
||||
#include "cpuset.h"
|
||||
#include <sched.h>
|
||||
#include <fcntl.h>
|
||||
#include <unistd.h>
|
||||
@@ -65,18 +64,18 @@ bool cmdOptionExists(char** begin, char** end, const std::string& option) {
|
||||
|
||||
const char *model_descriptions[] = {
|
||||
"single node VEGA20 4P1H",
|
||||
"single node VEGA20 4P1H Alt. Model",
|
||||
"single node VEGA20 4P2H",
|
||||
"single node gfx908 4P3L",
|
||||
"single node gfx908 8P6L",
|
||||
"single node gfx908 8P6L Alt. Connection",
|
||||
"single node 8 VEGA20 PCIe on Rome",
|
||||
"single node gfx908 8P6L on Rome",
|
||||
"single node 8 VEGA20 PCIe",
|
||||
"4 nodes with 8 GPUs PCIe 1 NIC",
|
||||
"4 nodes with 8 GPUs PCIe 1 NIC 2nd PLX Bridge",
|
||||
"4 nodes with 8 GPUs PCIe 2 NIC",
|
||||
"2 nodes VEGA20 4P1H",
|
||||
"4 nodes with 8 VEGA20 GPUs XGMI 4P2H 1 NIC",
|
||||
"4 nodes 8 GPUs PCIe 2 NICs on Rome",
|
||||
"3 nodes 8 GPUs PCIe + 1 Rome 8 GPUs PCIe + 2 nodes gfx908 4P3L",
|
||||
"4 nodes with 8 VEGA20 GPUs XGMI 4P2H 1 NIC 2nd Hive",
|
||||
"4 nodes with 8 VEGA20 GPUs XGMI 4P2H 2 NIC",
|
||||
NULL,
|
||||
};
|
||||
|
||||
@@ -97,97 +96,75 @@ int main(int argc,char* argv[])
|
||||
if (mi)
|
||||
model_id = atol(mi);
|
||||
|
||||
// CPU, GPU and NIC devices on Skylake
|
||||
CpuDevices skylake("Skylake", SKL_QPI_WIDTH, SKL_CPUPCI_WIDTH, SKL_PCI_WIDTH);
|
||||
GpuDevices vg20_pcie(8, busIds_8, gpuPciPaths_8, gpuPciNumaIds_8, conn_mat_pcie);
|
||||
GpuDevices vg20_4p1h(4, busIds_8, gpuPciPaths_8, gpuPciNumaIds_8, conn_mat_4p2h);
|
||||
GpuDevices vg20_4p2h(8, busIds_8, gpuPciPaths_8, gpuPciNumaIds_8, conn_mat_4p2h);
|
||||
GpuDevices gfx908_4p3l(4, busIds_8, gpuPciPaths_8, gpuPciNumaIds_8, conn_mat_8p6l);
|
||||
GpuDevices gfx908_8p6l(8, busIds_8, gpuPciPaths_8, gpuPciNumaIds_8, conn_mat_8p6l);
|
||||
GpuDevices gfx908_8p6l_1(8, busIds_8, gpuPciPaths_8, gpuPciNumaIds_8, conn_mat_8p6l_1);
|
||||
NetDevices nic_1(1, netPciPaths_1, netGuids_1, netPciNumaIds_1);
|
||||
NetDevices nic_1_1(1, netPciPaths_1_1, netGuids_1, netPciNumaIds_1);
|
||||
NetDevices nic_2(2, netPciPaths_2, netGuids_2, netPciNumaIds_2);
|
||||
|
||||
// CPU, GPU and NIC devices on Rome
|
||||
CpuDevices rome("Rome", ROME_QPI_WIDTH, ROME_CPUPCI_WIDTH, ROME_PCI_WIDTH);
|
||||
GpuDevices vg20_pcie_rome(8, rome_busIds_8, rome_gpuPciPaths_8, rome_gpuPciNumaIds_8, conn_mat_rome);
|
||||
NetDevices nic_1_rome(1, rome_netPciPaths_1, rome_netGuids_1, rom_netPciNumaIds_1);
|
||||
NetDevices nic_2_rome(2, rome_netPciPaths_2, rome_netGuids_2, rom_netPciNumaIds_2);
|
||||
|
||||
// 8 GPUs PCIe 1 NIC
|
||||
NodeModel model_8pcie_1nic(skylake, vg20_pcie, nic_1, "Skylake 8 GPUs PCIe");
|
||||
|
||||
// 8 GPUs PCIe 2 NIC
|
||||
NodeModel model_8pcie_2nic(skylake, vg20_pcie, nic_2, "Skylake 8 GPUs PCIe 2 NIC");
|
||||
|
||||
// VEGA20 4P1H, use VEGA20 4P2H model
|
||||
NodeModel model_vg20_4p1h_1nic(skylake, vg20_4p1h, nic_1, "Skylake VEGA20 4P1H");
|
||||
|
||||
// VEGA20 GPUs XGMI 4P2H 1 NIC
|
||||
NodeModel model_vg20_4p2h_1nic(skylake, vg20_4p2h, nic_1_1, "Skylake VEGA20 4P2H");
|
||||
|
||||
// gfx908 4P3L
|
||||
NodeModel model_gfx908_4p_1nic(skylake, gfx908_4p3l, nic_1, "Skylake gfx908 4P3L");
|
||||
|
||||
// gfx908 8P6L
|
||||
NodeModel model_gfx908_8p_1nic(skylake, gfx908_8p6l, nic_1, "Skylake gfx908 8P6L");
|
||||
|
||||
// gfx908 8P6L alternative connection
|
||||
NodeModel model_gfx908_8p_1nic_1(skylake, gfx908_8p6l_1, nic_1, "Skylake gfx908 8P6L Alt. Connection");
|
||||
|
||||
// 8 GPUs PCIe on Rome
|
||||
NodeModel model_8pcie_1nic_rome(rome, vg20_pcie_rome, nic_1_rome, "Rome 8 GPUs PCIe");
|
||||
|
||||
// 8 GPUs PCIe 2 NICs on Rome
|
||||
NodeModel model_8pcie_2nic_rome(rome, vg20_pcie_rome, nic_2_rome, "Rome 8 GPUs PCIe 2 NICs");
|
||||
|
||||
// gfx908 8P6L on Rome
|
||||
NodeModel model_gfx908_8p_1nic_rome(rome, gfx908_8p6l, nic_1, "Rome gfx908 8P6L");
|
||||
|
||||
NetworkModel network;
|
||||
NodeModel* node;
|
||||
|
||||
switch(model_id) {
|
||||
case 0:
|
||||
network.AddNode(model_vg20_4p1h_1nic);
|
||||
node = new NodeModel("topo_4p1h.xml");
|
||||
network.AddNode(node);
|
||||
break;
|
||||
case 1:
|
||||
network.AddNode(model_vg20_4p2h_1nic);
|
||||
node = new NodeModel("topo_4p1h_1.xml");
|
||||
network.AddNode(node);
|
||||
break;
|
||||
case 2:
|
||||
network.AddNode(model_gfx908_4p_1nic);
|
||||
node = new NodeModel("topo_4p2h.xml");
|
||||
network.AddNode(node);
|
||||
break;
|
||||
case 3:
|
||||
network.AddNode(model_gfx908_8p_1nic);
|
||||
node = new NodeModel("topo_4p3l.xml");
|
||||
network.AddNode(node);
|
||||
break;
|
||||
case 4:
|
||||
network.AddNode(model_gfx908_8p_1nic_1);
|
||||
node = new NodeModel("topo_8p6l.xml");
|
||||
network.AddNode(node);
|
||||
break;
|
||||
case 5:
|
||||
network.AddNode(model_8pcie_1nic_rome);
|
||||
node = new NodeModel("topo_8p_pcie.xml");
|
||||
network.AddNode(node);
|
||||
break;
|
||||
case 6:
|
||||
network.AddNode(model_gfx908_8p_1nic_rome);
|
||||
for (int i=0; i<4; i++) {
|
||||
node = new NodeModel("topo_8p_pcie.xml");
|
||||
network.AddNode(node);
|
||||
}
|
||||
break;
|
||||
case 7:
|
||||
for (int i = 0; i < 4; i ++) network.AddNode(model_8pcie_1nic);
|
||||
for (int i=0; i<4; i++) {
|
||||
node = new NodeModel("topo_8p_pcie_1.xml");
|
||||
network.AddNode(node);
|
||||
}
|
||||
break;
|
||||
case 8:
|
||||
for (int i = 0; i < 4; i ++) network.AddNode(model_8pcie_2nic);
|
||||
for (int i=0; i<4; i++) {
|
||||
node = new NodeModel("topo_8p_pcie_2nic.xml");
|
||||
network.AddNode(node);
|
||||
}
|
||||
break;
|
||||
case 9:
|
||||
for (int i = 0; i < 2; i ++) network.AddNode(model_vg20_4p1h_1nic);
|
||||
for (int i=0; i<2; i++) {
|
||||
node = new NodeModel("topo_4p1h.xml");
|
||||
network.AddNode(node);
|
||||
}
|
||||
break;
|
||||
case 10:
|
||||
for (int i = 0; i < 4; i ++) network.AddNode(model_vg20_4p2h_1nic);
|
||||
for (int i=0; i<4; i++) {
|
||||
node = new NodeModel("topo_4p2h.xml");
|
||||
network.AddNode(node);
|
||||
}
|
||||
break;
|
||||
case 11:
|
||||
for (int i = 0; i < 4; i ++) network.AddNode(model_8pcie_2nic_rome);
|
||||
for (int i=0; i<4; i++) {
|
||||
node = new NodeModel("topo_4p2h_1.xml");
|
||||
network.AddNode(node);
|
||||
}
|
||||
break;
|
||||
case 12:
|
||||
for (int i = 0; i < 3; i ++) network.AddNode(model_8pcie_1nic);
|
||||
network.AddNode(model_8pcie_1nic_rome);
|
||||
for (int i = 0; i < 2; i ++) network.AddNode(model_gfx908_4p_1nic);
|
||||
for (int i=0; i<4; i++) {
|
||||
node = new NodeModel("topo_4p2h_2nic.xml");
|
||||
network.AddNode(node);
|
||||
}
|
||||
break;
|
||||
default:
|
||||
printf("Invalid model_id %d\n", model_id);
|
||||
@@ -203,8 +180,8 @@ int main(int argc,char* argv[])
|
||||
for (int i = 0; i < nranks; i++) {
|
||||
node_model = network.GetNode(i);
|
||||
assert(node_model!=0);
|
||||
printf("Rank %d: node %d (%s) GPU busId %lx\n", i, node_model->nodeId,
|
||||
node_model->description, node_model->getGpuBusId(node_model->rankToCudaDev(i)));
|
||||
printf("Rank %d: node %d cudaDev %d GPU busId %lx\n", i, node_model->nodeId,
|
||||
node_model->rankToCudaDev(i), node_model->getGpuBusId(i));
|
||||
}
|
||||
|
||||
NCCLCHECK(ncclCalloc(&comm, nranks));
|
||||
@@ -220,21 +197,22 @@ int main(int argc,char* argv[])
|
||||
comm[i].nRanks = nranks;
|
||||
node_model = network.GetNode(i);
|
||||
assert(node_model!=0);
|
||||
comm[i].topo = node_model->getSystem(i);
|
||||
bootstrapAllGather(&comm[i], allGather1Data);
|
||||
}
|
||||
|
||||
struct ncclTopoGraph treeGraph, ringGraph;
|
||||
struct ncclTopoGraph treeGraph, ringGraph, collNetGraph;
|
||||
|
||||
for (int i = 0; i < nranks; i++) {
|
||||
node_model = network.GetNode(i);
|
||||
assert(node_model!=0);
|
||||
initTransportsRank_1(&comm[i], allGather1Data, allGather3Data, treeGraph, ringGraph);
|
||||
initTransportsRank_1(&comm[i], allGather1Data, allGather3Data, treeGraph, ringGraph, collNetGraph);
|
||||
}
|
||||
|
||||
for (int i = 0; i < nranks; i++) {
|
||||
node_model = network.GetNode(i);
|
||||
assert(node_model!=0);
|
||||
initTransportsRank_3(&comm[i], allGather3Data, treeGraph, ringGraph);
|
||||
initTransportsRank_3(&comm[i], allGather3Data, treeGraph, ringGraph, collNetGraph);
|
||||
}
|
||||
|
||||
free(allGather3Data);
|
||||
|
||||
@@ -25,12 +25,19 @@
|
||||
#include <sys/types.h>
|
||||
#include <sys/stat.h>
|
||||
#include <unistd.h>
|
||||
#include "xml.h"
|
||||
#include "coll_net.h"
|
||||
#include "model.h"
|
||||
#include "utils.h"
|
||||
|
||||
extern NodeModel *node_model;
|
||||
|
||||
NCCL_PARAM(CrossNic, "CROSS_NIC", 2);
|
||||
NCCL_PARAM(CollNetEnable, "COLLNET_ENABLE", 0);
|
||||
NCCL_PARAM(GraphDumpFileRank, "GRAPH_DUMP_FILE_RANK", 0);
|
||||
|
||||
thread_local int ncclDebugNoWarn = 0;
|
||||
ncclCollNet_t* ncclCollNet = NULL;
|
||||
|
||||
// Get current Compute Capability
|
||||
int ncclCudaCompCap() {
|
||||
@@ -43,7 +50,7 @@ ncclResult_t int64ToBusId(int64_t id, char* busId) {
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
ncclResult_t busIdToInt64(char* busId, int64_t* id) {
|
||||
ncclResult_t busIdToInt64(const char* busId, int64_t* id) {
|
||||
const int size = strlen(busId);
|
||||
char* hexStr;
|
||||
NCCLCHECK(ncclCalloc(&hexStr, size));
|
||||
@@ -87,9 +94,9 @@ void ncclDebugLog(ncclDebugLogLevel level, unsigned long flags, const char *file
|
||||
if (ncclDebugLevel == -1) ncclDebugInit();
|
||||
if (level == NCCL_LOG_TRACE && ncclDebugLevel != NCCL_LOG_TRACE) return;
|
||||
char buffer[1024];
|
||||
size_t len;
|
||||
len = snprintf(buffer, sizeof(buffer),
|
||||
"[%d:%d] ", node_model->nodeId, node_model->currRank);
|
||||
size_t len = 0;
|
||||
if (node_model) len = snprintf(buffer, sizeof(buffer),
|
||||
"[%d:%d] ", node_model->nodeId, node_model->currRank);
|
||||
va_list args;
|
||||
va_start(args, fmt);
|
||||
vsprintf(buffer+len, fmt, args);
|
||||
@@ -102,6 +109,16 @@ void ncclDebugLog(ncclDebugLogLevel level, unsigned long flags, const char *file
|
||||
}
|
||||
}
|
||||
|
||||
ncclResult_t ncclTopoGetSystem(const char* xmlTopoFile, struct ncclTopoSystem** system) {
|
||||
struct ncclXml* xml;
|
||||
NCCLCHECK(ncclCalloc(&xml, 1));
|
||||
NCCLCHECK(ncclTopoGetXmlFromFile(xmlTopoFile, xml));
|
||||
NCCLCHECK(ncclTopoGetSystemFromXml(xml, system));
|
||||
free(xml);
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
|
||||
ncclResult_t bootstrapAllGather(struct ncclComm* comm, struct allGather1Data_t * allGather1Data) {
|
||||
// AllGather1 - begin
|
||||
allGather1Data[comm->rank].peerInfo.rank = comm->rank;
|
||||
@@ -110,12 +127,12 @@ ncclResult_t bootstrapAllGather(struct ncclComm* comm, struct allGather1Data_t *
|
||||
allGather1Data[comm->rank].peerInfo.hostHash = node_model->hostHash;
|
||||
allGather1Data[comm->rank].peerInfo.pidHash = node_model->pidHash;
|
||||
allGather1Data[comm->rank].peerInfo.shmDev = 0x19;
|
||||
allGather1Data[comm->rank].peerInfo.busId = node_model->getGpuBusId(node_model->rankToCudaDev(comm->rank));
|
||||
allGather1Data[comm->rank].peerInfo.busId = node_model->getGpuBusId(comm->rank);
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather1Data_t *allGather1Data,
|
||||
struct allGather3Data_t *allGather3Data, struct ncclTopoGraph& treeGraph, struct ncclTopoGraph& ringGraph) {
|
||||
ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather1Data_t *allGather1Data, struct allGather3Data_t *allGather3Data,
|
||||
struct ncclTopoGraph& treeGraph, struct ncclTopoGraph& ringGraph, struct ncclTopoGraph& collNetGraph) {
|
||||
// We use 3 AllGathers
|
||||
// 1. { peerInfo, comm }
|
||||
// 2. ConnectTransport[nranks], ConnectValue[nranks]
|
||||
@@ -147,45 +164,70 @@ ncclResult_t initTransportsRank_1(struct ncclComm* comm, struct allGather1Data_t
|
||||
// AllGather1 - end
|
||||
|
||||
// Topo detection / System graph creation
|
||||
NCCLCHECK(ncclTopoGetSystem(comm, &comm->topo));
|
||||
//NCCLCHECK(ncclTopoGetSystem(comm, &comm->topo));
|
||||
// Compute paths between GPUs and NICs
|
||||
NCCLCHECK(ncclTopoComputePaths(comm->topo, comm->peerInfo));
|
||||
// Remove inaccessible GPUs and unused NICs
|
||||
NCCLCHECK(ncclTopoTrimSystem(comm->topo, comm));
|
||||
// Recompute paths after trimming
|
||||
NCCLCHECK(ncclTopoComputePaths(comm->topo, comm->peerInfo));
|
||||
// Compute max speed to accelerate search
|
||||
NCCLCHECK(ncclTopoGetMaxSpeed(comm->topo));
|
||||
// Init search
|
||||
NCCLCHECK(ncclTopoSearchInit(comm->topo));
|
||||
// Print final topology
|
||||
NCCLCHECK(ncclTopoPrint(comm->topo));
|
||||
|
||||
// Get rings and trees
|
||||
//struct ncclTopoGraph treeGraph;
|
||||
treeGraph.pattern = NCCL_TOPO_PATTERN_SPLIT_TREE;
|
||||
treeGraph.crossNic = ncclParamCrossNic();
|
||||
// We communicate only half the data between node with trees on 2 nodes.
|
||||
NCCLCHECK(ncclTopoCompute(comm->topo, &treeGraph));
|
||||
NCCLCHECK(ncclTopoPrintGraph(comm->topo, &treeGraph));
|
||||
//struct ncclTopoGraph ringGraph;
|
||||
ringGraph.id = 0;
|
||||
ringGraph.pattern = NCCL_TOPO_PATTERN_RING;
|
||||
ringGraph.crossNic = ncclParamCrossNic();
|
||||
ringGraph.collNet = 0;
|
||||
ringGraph.minChannels = 1;
|
||||
ringGraph.maxChannels = MAXCHANNELS/2;
|
||||
NCCLCHECK(ncclTopoCompute(comm->topo, &ringGraph));
|
||||
NCCLCHECK(ncclTopoPrintGraph(comm->topo, &ringGraph));
|
||||
|
||||
//struct ncclTopoGraph treeGraph;
|
||||
treeGraph.id = 1;
|
||||
treeGraph.pattern = NCCL_TOPO_PATTERN_SPLIT_TREE;
|
||||
treeGraph.crossNic = ncclParamCrossNic();
|
||||
treeGraph.collNet = 0;
|
||||
treeGraph.minChannels = comm->topo->nodes[NET].count != 0 ? 1 : ringGraph.nChannels;
|
||||
treeGraph.maxChannels = ringGraph.nChannels;
|
||||
NCCLCHECK(ncclTopoCompute(comm->topo, &treeGraph));
|
||||
NCCLCHECK(ncclTopoPrintGraph(comm->topo, &treeGraph));
|
||||
|
||||
//struct ncclTopoGraph collNetGraph;
|
||||
collNetGraph.id = 2;
|
||||
collNetGraph.pattern = NCCL_TOPO_PATTERN_TREE;
|
||||
collNetGraph.collNet = 1;
|
||||
collNetGraph.crossNic = ncclParamCrossNic();
|
||||
collNetGraph.minChannels = collNetGraph.maxChannels = ringGraph.nChannels;
|
||||
NCCLCHECK(ncclTopoCompute(comm->topo, &collNetGraph));
|
||||
NCCLCHECK(ncclTopoPrintGraph(comm->topo, &collNetGraph));
|
||||
|
||||
if (comm->rank == ncclParamGraphDumpFileRank()) {
|
||||
struct ncclTopoGraph* graphs[3] = { &ringGraph, &treeGraph, &collNetGraph };
|
||||
NCCLCHECK(ncclTopoDumpGraphs(comm->topo, 3, graphs));
|
||||
}
|
||||
|
||||
// AllGather3 - begin
|
||||
allGather3Data[rank].cudaCompCap = ncclCudaCompCap();
|
||||
allGather3Data[rank].nvlink = treeGraph.nvlink;
|
||||
allGather3Data[rank].nChannels = comm->nChannels = std::min(treeGraph.nChannels, ringGraph.nChannels);
|
||||
allGather3Data[rank].tree.sameChannels = treeGraph.sameChannels;
|
||||
allGather3Data[rank].tree.speedIntra = treeGraph.speedIntra;
|
||||
allGather3Data[rank].tree.speedInter = treeGraph.speedInter;
|
||||
allGather3Data[rank].tree.nvlink = treeGraph.nvlink;
|
||||
allGather3Data[rank].tree.typeIntra = treeGraph.typeIntra;
|
||||
allGather3Data[rank].ring.sameChannels = ringGraph.sameChannels;
|
||||
allGather3Data[rank].ring.speedIntra = ringGraph.speedIntra;
|
||||
allGather3Data[rank].ring.speedInter = ringGraph.speedInter;
|
||||
allGather3Data[rank].ring.nvlink = ringGraph.nvlink;
|
||||
allGather3Data[rank].ring.typeIntra = ringGraph.typeIntra;
|
||||
allGather3Data[rank].collNet.sameChannels = collNetGraph.sameChannels;
|
||||
allGather3Data[rank].collNet.speedIntra = collNetGraph.speedIntra;
|
||||
allGather3Data[rank].collNet.speedInter = collNetGraph.speedInter;
|
||||
allGather3Data[rank].collNet.typeIntra = collNetGraph.typeIntra;
|
||||
|
||||
NCCLCHECK(ncclTopoPreset(comm, &treeGraph, &ringGraph, &allGather3Data[rank].topoRanks));
|
||||
NCCLCHECK(ncclTopoPreset(comm, &treeGraph, &ringGraph, &collNetGraph, &allGather3Data[rank].topoRanks));
|
||||
//INFO(NCCL_GRAPH, "%d: nvlink %d nChannels %d tree.sameChannels %d tree.speedIntra %d tree.speedInter %d tree.nvlink %d ring.sameChannels %d ring.speedIntra %d ring.speedInter %d ring.nvlink %d",
|
||||
// rank, allGather3Data[rank].nvlink, allGather3Data[rank].nChannels, allGather3Data[rank].tree.sameChannels, allGather3Data[rank].tree.speedIntra, allGather3Data[rank].tree.speedInter, allGather3Data[rank].tree.nvlink,
|
||||
// allGather3Data[rank].ring.sameChannels, allGather3Data[rank].ring.speedIntra, allGather3Data[rank].ring.speedInter, allGather3Data[rank].ring.nvlink);
|
||||
@@ -203,14 +245,8 @@ static ncclResult_t selectTransport(struct ncclTopoSystem* topo, struct ncclTopo
|
||||
int ret = 0;
|
||||
NCCLCHECK(transport->canConnect(&ret, topo, graph, myInfo, peerInfo));
|
||||
if (ret) {
|
||||
//cpu_set_t affinitySave;
|
||||
//sched_getaffinity(0, sizeof(cpu_set_t), &affinitySave);
|
||||
//int cudaDev;
|
||||
//CUDACHECK(hipGetDevice(&cudaDev));
|
||||
//setCpuAffinity(cudaDev);
|
||||
connector->transportComm = transportComm;
|
||||
NCCLCHECK(transportComm->setup(topo, graph, myInfo, peerInfo, connect, connector, buffSize, channelId));
|
||||
//sched_setaffinity(0, sizeof(cpu_set_t), &affinitySave);
|
||||
return ncclSuccess;
|
||||
}
|
||||
}
|
||||
@@ -265,21 +301,26 @@ static ncclResult_t p2pSetup(struct ncclComm* comm, struct ncclTopoGraph* graph,
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
|
||||
ncclResult_t initChannel(struct ncclComm* comm, int channelid) {
|
||||
struct ncclChannel* channel = comm->channels+channelid;
|
||||
channel->id = channelid;
|
||||
|
||||
// Setup intermediate buffering
|
||||
//channel->buffSize = ncclParamBuffsize();
|
||||
//int buffSize = ncclParamBuffsize();
|
||||
int cpuArch, cpuVendor, cpuModel;
|
||||
NCCLCHECK(ncclTopoCpuType(comm->topo, &cpuArch, &cpuVendor, &cpuModel));
|
||||
//channel->buffSize = buffSize != -2 ? buffSize :
|
||||
// cpuArch == NCCL_TOPO_CPU_ARCH_ARM ? DEFAULT_BUFFER_SIZE_BYTES_ARM : DEFAULT_BUFFER_SIZE_BYTES;
|
||||
|
||||
// Ring index to user rank table.
|
||||
//NCCLCHECK(ncclCudaCalloc(&channel->ring.devUserRanks, comm->nRanks));
|
||||
NCCLCHECK(ncclCalloc(&channel->ring.userRanks, comm->nRanks));
|
||||
|
||||
// Communication structures with peers.
|
||||
//NCCLCHECK(ncclCudaCalloc(&channel->devPeers, comm->nRanks));
|
||||
NCCLCHECK(ncclCalloc(&channel->peers, comm->nRanks));
|
||||
for (size_t i=0; i<comm->nRanks; ++i) {
|
||||
//NCCLCHECK(ncclCudaCalloc(&channel->devPeers, comm->nRanks+1)); // The extra one rank is for collnet root (i.e. network)
|
||||
NCCLCHECK(ncclCalloc(&channel->peers, comm->nRanks+1));
|
||||
for (size_t i=0; i<comm->nRanks+1; ++i) {
|
||||
channel->peers[i].send.comm = comm;
|
||||
channel->peers[i].recv.comm = comm;
|
||||
}
|
||||
@@ -307,8 +348,8 @@ static ncclResult_t setupChannel(struct ncclComm* comm, int channelId, int rank,
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t *allGather3Data, struct ncclTopoGraph& treeGraph,
|
||||
struct ncclTopoGraph& ringGraph) {
|
||||
ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t *allGather3Data,
|
||||
struct ncclTopoGraph& treeGraph, struct ncclTopoGraph& ringGraph, struct ncclTopoGraph& collNetGraph) {
|
||||
int rank = comm->rank;
|
||||
int nranks = comm->nRanks;
|
||||
//NCCLCHECK(bootstrapAllGather(comm->bootstrap, allGather3Data, sizeof(*allGather3Data)));
|
||||
@@ -329,6 +370,15 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t
|
||||
if (i == comm->rank) comm->node = node;
|
||||
}
|
||||
|
||||
char line[1024];
|
||||
sprintf(line, "nodesFirstRank: ");
|
||||
int offset = strlen(line);
|
||||
for (int i=0; i<comm->nNodes; i++) {
|
||||
sprintf(line+offset, "%d ", nodesFirstRank[i]);
|
||||
offset = strlen(line);
|
||||
}
|
||||
INFO(NCCL_INIT, "%s", line);
|
||||
|
||||
// Determine the minimum CUDA Compute capability of all GPUs
|
||||
int myCompCap = allGather3Data[rank].cudaCompCap;
|
||||
int minCompCap = myCompCap, maxCompCap = myCompCap;
|
||||
@@ -337,9 +387,6 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t
|
||||
maxCompCap = std::max(allGather3Data[i].cudaCompCap, maxCompCap);
|
||||
}
|
||||
|
||||
comm->nvlink = 1;
|
||||
for (int i = 0; i < nranks; i++) comm->nvlink &= allGather3Data[i].nvlink;
|
||||
|
||||
int nChannelsOrig = comm->nChannels;
|
||||
struct ncclTopoRanks** allTopoRanks;
|
||||
NCCLCHECK(ncclCalloc(&allTopoRanks, comm->nRanks));
|
||||
@@ -350,11 +397,15 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t
|
||||
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.nvlink = std::min(allGather3Data[i].tree.nvlink, treeGraph.nvlink);
|
||||
treeGraph.typeIntra = std::min(allGather3Data[i].tree.typeIntra, treeGraph.typeIntra);
|
||||
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.nvlink = std::min(allGather3Data[i].ring.nvlink, ringGraph.nvlink);
|
||||
ringGraph.typeIntra = std::min(allGather3Data[i].ring.typeIntra, ringGraph.typeIntra);
|
||||
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.typeIntra = std::min(allGather3Data[i].collNet.typeIntra, collNetGraph.typeIntra);
|
||||
}
|
||||
|
||||
if (comm->nChannels < nChannelsOrig) {
|
||||
@@ -366,24 +417,23 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t
|
||||
int *rings;
|
||||
NCCLCHECK(ncclCalloc(&rings, nranks*MAXCHANNELS));
|
||||
|
||||
char line[1024];
|
||||
sprintf(line, "nodesFirstRank: ");
|
||||
int offset = strlen(line);
|
||||
for (int i=0; i<comm->nNodes; i++) {
|
||||
sprintf(line+offset, "%d ", nodesFirstRank[i]);
|
||||
offset = strlen(line);
|
||||
}
|
||||
INFO(NCCL_INIT, "%s", line);
|
||||
|
||||
NCCLCHECK(ncclTopoPostset(comm, nodesFirstRank, allTopoRanks, rings));
|
||||
if (comm->nNodes > 1 &&
|
||||
ncclParamCollNetEnable() == 1 &&
|
||||
collNetSupport()) {
|
||||
NCCLCHECK(ncclTopoConnectCollNet(comm, &collNetGraph, rank));
|
||||
}
|
||||
|
||||
free(allTopoRanks);
|
||||
free(nodesFirstRank);
|
||||
//free(allGather3Data);
|
||||
|
||||
// AllGather3 - end
|
||||
|
||||
TRACE(NCCL_INIT, "rank %d nranks %d - BUILT %d TREES/RINGS", rank, nranks, comm->nChannels);
|
||||
|
||||
NCCLCHECK(ncclTopoSetThresholds(comm, minCompCap, maxCompCap, &treeGraph, &ringGraph, &collNetGraph));
|
||||
|
||||
line[0]='\0';
|
||||
for (int c=0; c<comm->nChannels; c++) {
|
||||
struct ncclTree* treeUp = &comm->channels[c].treeUp;
|
||||
@@ -395,26 +445,56 @@ ncclResult_t initTransportsRank_3(struct ncclComm* comm, struct allGather3Data_t
|
||||
line[1023] = '\0';
|
||||
INFO(NCCL_INIT, "Trees%s", line);
|
||||
|
||||
free(rings);
|
||||
|
||||
// Done with AllGather1 data
|
||||
//free(allGather1Data);
|
||||
|
||||
TRACE(NCCL_INIT, "rank %d nranks %d - DONE", rank, nranks);
|
||||
// Set Affinity to a CPU local the our GPU, so that all memory we allocate
|
||||
// on the host is local.
|
||||
cpu_set_t affinitySave;
|
||||
sched_getaffinity(0, sizeof(cpu_set_t), &affinitySave);
|
||||
NCCLCHECK(ncclTopoSetAffinity(comm->topo, comm->rank));
|
||||
ncclResult_t ret;
|
||||
|
||||
// Connect with prev/next for each ring
|
||||
struct ncclConnect *connect;
|
||||
NCCLCHECK(ncclCalloc(&connect, 2));
|
||||
NCCLCHECKGOTO(ncclCalloc(&connect, 2), ret, affinity_restore);
|
||||
for (int c=0; c<comm->nChannels; c++) {
|
||||
struct ncclChannel* channel = comm->channels+c;
|
||||
NCCLCHECK(setupChannel(comm, c, rank, nranks, rings+c*nranks));
|
||||
NCCLCHECKGOTO(setupChannel(comm, c, rank, nranks, rings+c*nranks), ret, affinity_restore);
|
||||
if (comm->nRanks == 1) continue;
|
||||
NCCLCHECK(p2pSetup(comm, &ringGraph, channel, 1, &channel->ring.prev, 1, &channel->ring.next));
|
||||
NCCLCHECK(p2pSetup(comm, &treeGraph, channel, NCCL_MAX_TREE_ARITY, channel->treeUp.down, 1, &channel->treeUp.up));
|
||||
NCCLCHECK(p2pSetup(comm, &treeGraph, channel, 1, &channel->treeDn.up, NCCL_MAX_TREE_ARITY, channel->treeDn.down));
|
||||
NCCLCHECKGOTO(p2pSetup(comm, &ringGraph, channel, 1, &channel->ring.prev, 1, &channel->ring.next), ret, affinity_restore);
|
||||
NCCLCHECKGOTO(p2pSetup(comm, &treeGraph, channel, NCCL_MAX_TREE_ARITY, channel->treeUp.down, 1, &channel->treeUp.up), ret, affinity_restore);
|
||||
NCCLCHECKGOTO(p2pSetup(comm, &treeGraph, channel, 1, &channel->treeDn.up, NCCL_MAX_TREE_ARITY, channel->treeDn.down), ret, affinity_restore);
|
||||
}
|
||||
|
||||
// Check if we can setup CollNet
|
||||
#if 0
|
||||
if (comm->nNodes > 1 &&
|
||||
ncclParamCollNetEnable() == 1 &&
|
||||
collNetSupport()) {
|
||||
int logicChannels = comm->nChannels/2;
|
||||
int collNetSetupFail = 0;
|
||||
const int recvIndex = 0; // recv GPU index is always 0
|
||||
const int sendIndex = collNetGraph.pattern == NCCL_TOPO_PATTERN_TREE ? 0 : 1; // send GPU index depends on topo pattern
|
||||
for (int c=0; c<logicChannels; c++) {
|
||||
struct ncclChannel* channelRecv = comm->channels+logicChannels+c;
|
||||
struct ncclChannel* channelSend = comm->channels+c;
|
||||
NCCLCHECK(p2pSetup(comm, &collNetGraph, channelRecv, 1, &channelRecv->collTreeDn.up, 1, channelRecv->collTreeDn.down));
|
||||
NCCLCHECK(p2pSetup(comm, &collNetGraph, channelSend, 1, channelSend->collTreeUp.down, 1, &channelSend->collTreeUp.up));
|
||||
const int recvMaster = collNetGraph.intra[c*comm->localRanks+recvIndex];
|
||||
const int sendMaster = collNetGraph.intra[c*comm->localRanks+sendIndex];
|
||||
if (collNetSetup(comm, &collNetGraph, channelRecv, logicChannels, rank, nranks, recvMaster, sendMaster, comm->nNodes, 1) != 1)
|
||||
collNetSetupFail = 1;
|
||||
if (collNetSetup(comm, &collNetGraph, channelSend, logicChannels, rank, nranks, sendMaster, recvMaster, comm->nNodes, 0) != 1)
|
||||
collNetSetupFail = 1;
|
||||
}
|
||||
// Verify CollNet setup across ranks
|
||||
NCCLCHECK(checkCollNetSetup(comm, rank, collNetSetupFail));
|
||||
}
|
||||
#endif
|
||||
TRACE(NCCL_INIT, "rank %d nranks %d - CONNECTED %d RINGS AND TREES", rank, nranks, comm->nChannels);
|
||||
free(connect);
|
||||
free(rings);
|
||||
|
||||
affinity_restore:
|
||||
if (ret != ncclSuccess) return ret;
|
||||
|
||||
return ncclSuccess;
|
||||
}
|
||||
|
||||
Ссылка в новой задаче
Block a user