diff options
Diffstat (limited to 'src/graph/topo.h')
-rw-r--r-- | src/graph/topo.h | 146 |
1 files changed, 76 insertions, 70 deletions
diff --git a/src/graph/topo.h b/src/graph/topo.h index 6b8a2f9..848fc03 100644 --- a/src/graph/topo.h +++ b/src/graph/topo.h @@ -1,5 +1,5 @@ /************************************************************************* - * Copyright (c) 2016-2019, NVIDIA CORPORATION. All rights reserved. + * Copyright (c) 2016-2020, NVIDIA CORPORATION. All rights reserved. * * See LICENSE.txt for license information ************************************************************************/ @@ -9,22 +9,24 @@ #include "graph.h" #include "core.h" - -#define LOC_WIDTH 5000 -#define PASCAL_NVLINK_WIDTH 18 -#define VOLTA_NVLINK_WIDTH 21 -#define PCI_WIDTH 12 // PCI Gen3 x16 -#define QPI_WIDTH 8 -#define SKL_QPI_WIDTH 12 -#define P9_WIDTH 32 -#define NET_WIDTH 12 // 100Gbit - -// Intel CPU convert GPU P2P traffic into 64B PCI TLPs, to GPU -// to GPU traffic consumed more PCI bandwidth. +#include <sched.h> + +#define LOC_WIDTH 5000.0 +#define PASCAL_NVLINK_WIDTH 18.0 +#define VOLTA_NVLINK_WIDTH 21.0 +#define PCI_WIDTH 12.0 // PCI Gen3 x16 +#define QPI_WIDTH 6.0 +#define SKL_QPI_WIDTH 9.0 +#define P9_WIDTH 32.0 +#define ARM_WIDTH 6.0 +#define NET_WIDTH 12.0 // 100Gbit + +// Intel CPU convert GPU P2P traffic into 64B PCI TLPs, so GPU +// to GPU traffic consumes more PCI bandwidth. #define INTEL_P2P(speed) (speed*9/12) #define INTEL_P2P_OVERHEAD(speed) (speed*12/9) -#define NCCL_TOPO_NODE_TYPES 6 +#define NCCL_TOPO_NODE_TYPES 7 #define GPU 0 #define PCI 1 #define NVS 2 @@ -33,37 +35,72 @@ #define NET 5 extern const char* topoNodeTypeStr[]; +// We want link types and path types to match as much as possible #define LINK_LOC 0 #define LINK_NVL 1 #define LINK_PCI 2 -#define LINK_QPI 3 -#define LINK_NET 4 +// Skipping 3 for PATH_PXB +// Skipping 4 for PATH_PHB +#define LINK_SYS 5 +#define LINK_NET 6 extern const char* topoLinkTypeStr[]; +#define PATH_LOC 0 +#define PATH_NVL 1 +#define PATH_PIX 2 +#define PATH_PXB 3 +#define PATH_PHB 4 +#define PATH_SYS 5 +#define PATH_NET 6 +extern const char* topoPathTypeStr[]; + struct ncclTopoNode; struct ncclTopoLink { int type; - int width; + float width; struct ncclTopoNode* remNode; }; #define NCCL_TOPO_MAX_LINKS 32 #define NCCL_TOPO_MAX_HOPS (NCCL_TOPO_MAX_NODES*NCCL_TOPO_NODE_TYPES) -#define SELECT_PATH 1 -#define SELECT_LAST 2 - -#define NET_GDR_MASK 0x70000000 struct ncclTopoLinkList { struct ncclTopoLink* list[NCCL_TOPO_MAX_HOPS]; int count; - int width; + float width; int type; }; +#define NCCL_TOPO_CPU_INTEL_BDW 1 +#define NCCL_TOPO_CPU_INTEL_SKL 2 + +#define NCCL_TOPO_UNDEF (-1) + struct ncclTopoNode { int type; int64_t id; - int rank; + // Type specific data + union { + struct { + int dev; // NVML dev number + int rank; + int cudaCompCap; + int gdrSupport; + }gpu; + struct { + uint64_t asic; + int port; + float width; + int gdrSupport; + int collSupport; + int maxChannels; + }net; + struct { + int arch; + int vendor; + int model; + cpu_set_t affinity; + }cpu; + }; int nlinks; struct ncclTopoLink links[NCCL_TOPO_MAX_LINKS]; // Pre-computed paths to GPUs and NICs @@ -79,60 +116,29 @@ struct ncclTopoNodeSet { struct ncclTopoSystem { struct ncclTopoNodeSet nodes[NCCL_TOPO_NODE_TYPES]; - int maxSpeed; - int maxWidth; - int searchInitDone; + float maxWidth; }; -static ncclResult_t ncclTopoCreateNode(struct ncclTopoSystem* system, struct ncclTopoNode** node, int type, uint64_t id) { +ncclResult_t ncclTopoGetNode(struct ncclTopoSystem* system, struct ncclTopoNode** node, int type, uint64_t id); +ncclResult_t ncclTopoCreateNode(struct ncclTopoSystem* system, struct ncclTopoNode** node, int type, uint64_t id); +ncclResult_t ncclTopoRemoveNode(struct ncclTopoSystem* system, int type, int id); +ncclResult_t ncclTopoConnectNodes(struct ncclTopoNode* node, struct ncclTopoNode* remNode, int type, float width); +ncclResult_t ncclTopoPrintPaths(struct ncclTopoSystem* system); +ncclResult_t ncclTopoLoadSystem(const char* xmlTopoFile, struct ncclTopoSystem* system); + +ncclResult_t ncclTopoGetSystemFromXml(struct ncclXml* xml, struct ncclTopoSystem** topoSystem); +ncclResult_t ncclTopoGetGraphFromXml(struct ncclXmlNode *xmlGraphs, struct ncclTopoSystem* system, struct ncclTopoGraph* graph); +ncclResult_t ncclTopoGetXmlFromGraphs(int ngraphs, struct ncclTopoGraph** graphs, struct ncclTopoSystem* system, struct ncclXml *xml); + +static ncclResult_t ncclTopoIdToIndex(struct ncclTopoSystem* system, int type, int64_t id, int* index) { + *index = -1; for (int i=0; i<system->nodes[type].count; i++) { if (system->nodes[type].nodes[i].id == id) { - *node = system->nodes[type].nodes+i; + *index = i; return ncclSuccess; } } - if (system->nodes[type].count == NCCL_TOPO_MAX_NODES) { - WARN("Error : tried to create too many nodes of type %d\n", type); - return ncclInternalError; - } - struct ncclTopoNode* n = system->nodes[type].nodes+system->nodes[type].count; - system->nodes[type].count++; - n->type = type; - n->id = id; - if (type == GPU) { - // Create link to itself (used in some corner cases) - n->nlinks=1; - n->links[0].type = LINK_LOC; - n->links[0].remNode = n; - n->links[0].width = LOC_WIDTH; - } - *node = n; - return ncclSuccess; -} - -static ncclResult_t ncclTopoConnectNodes(struct ncclTopoNode* node, struct ncclTopoNode* remNode, int type, int width) { - // Aggregate links into higher width for NVLink - struct ncclTopoLink* link; - for (link = node->links; link->remNode; link++) { - if (link->remNode == remNode && link->type == type) break; - } - if (link->remNode == NULL) node->nlinks++; - link->type = type; - link->remNode = remNode; - link->width += width; - - // Sort links in BW descending order - struct ncclTopoLink linkSave; - memcpy(&linkSave, link, sizeof(struct ncclTopoLink)); - while (link != node->links) { - if ((link-1)->width >= linkSave.width) break; - memcpy(link, link-1, sizeof(struct ncclTopoLink)); - link--; - } - memcpy(link, &linkSave, sizeof(struct ncclTopoLink)); - return ncclSuccess; + return ncclInternalError; } -ncclResult_t ncclTopoPrintPaths(struct ncclTopoSystem* system); - #endif |