NCCL拓扑管理 - Search模块

已于 2024-07-29 15:20:14 修改
阅读量1.8k 收藏 16
点赞数 32
分类专栏： NCCL源码剖析 文章标签： 人工智能
于 2024-07-19 09:12:29 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/weixin_42371021/article/details/140458256
版权

NCCL Search模块从节点连接和节点路径构造完毕的系统拓扑中搜索出各类逻辑拓扑（如环拓扑，树拓扑等）的通道（每个拓扑可能有多个通道，通道间可以并行运行集合通信算法）和节点（GPU节点和NET节点）。后续的connect模块将这些搜索出来的通道和节点根据拓扑的类型构建起来。

数据结构

ncclTopoGraph是search模块的主要结构体，它包含搜索的输入超参，比如id，pattern等，也包括了搜索结果的输出，比如通道数nChannels，每个通道的节点（非网络相连的节点信息保存在intra数组里，网络相连的节点保存在inter数组里）。
其中pattern包含以下5种。

#define NCCL_TOPO_PATTERN_BALANCED_TREE 1   // Spread NIC traffic between two GPUs (Tree parent + one child on first GPU, second child on second GPU)
#define NCCL_TOPO_PATTERN_SPLIT_TREE 2      // Spread NIC traffic between two GPUs (Tree parent on first GPU, tree children on the second GPU)
#define NCCL_TOPO_PATTERN_TREE 3            // All NIC traffic going to/from the same GPU
#define NCCL_TOPO_PATTERN_RING 4            // Ring
#define NCCL_TOPO_PATTERN_NVLS 5            // NVLS+SHARP and NVLS+Tree

ncclTopoGraph的成员介绍参见注释。

struct ncclTopoGraph {
  // Input / output of the search
  int id;                                             // ring : 0, tree : 1, collnet : 2
  int pattern;                                        // topo pattern
  int crossNic;                                       // if inbound Net and outbound Net use different NICs or not
  int collNet;										  // is CollNet or not
  int minChannels;									  // minChannel for search
  int maxChannels;									  // maxChannel for search
  
  // Output of the search
  int nChannels;									  // number of channels searched, can run in parallel
  float bwIntra;									  // intra-connect bandwidth
  float bwInter;									  // inter-connect bandwidth
  float latencyInter;								  // inter-connect latency
  int typeIntra; 									  // intra-connect type
  int typeInter; 									  // inter-connect type
  int sameChannels;                                   // 1 if all channels are the same type and bandwidth, 0 otherwise
  int nHops;                                          // Number of hops in the graph
  int intra[MAXCHANNELS*NCCL_TOPO_MAX_NODES];         // intra[channel_id][node_id] 
  int64_t inter[MAXCHANNELS*2];                       // inter[channel_id][2]，one for inbound and one for outbound
};

接口

ncclTopoSearchInit

ncclTopoSearchInit接口负责搜索前完成系统拓扑maxBw和totalBw的初始化，这两个数据也作为搜索的参考值。

// Setup system->maxBw and system->totalBw
// maxBw: maximum bandwidth between GPU and GPU or GPU and NET if cross NIC is required
// totalBw: maximum bandwidth a GPU node can have(NVLink or PCI)
ncclResult_t ncclTopoSearchInit(struct ncclTopoSystem* system) {
  system->maxBw = 0.0;
  system->totalBw = 0.0;
  int inter = system->nodes[NET].count;
  // If we have only one GPU, we can use local bandwidth
  if (inter == 0 && system->nodes[GPU].count == 1) {
    system->maxBw = LOC_BW;
    system->totalBw = LOC_BW;
    return ncclSuccess;
  }
  for (int g=0; g<system->nodes[GPU].count; g++) {
    struct ncclTopoNode* gpu = system->nodes[GPU].nodes+g;
    // Find the maximum bandwidth between GPU and GPU or GPU and NET
    system->maxBw = std::max(system->maxBw, getMaxBw(system, gpu, inter ? NET : GPU));
    system->totalBw = std::max(system->totalBw, getTotalBw(system, gpu));
  }
  return ncclSuccess;
}

ncclTopoCompute

ncclTopoCompute接口接收连接和路径构造完毕的ncclTopoSystem，搜索参数填写在ncclTopoGraph结构体中，将根据ncclTopoGraph中的预置信息展开搜索。
搜索速度变化的数值列举在以下数组中，数组下标越小速度越快。针对sm90有独立的速度数组。

float speedArrayIntra[] = { 40.0, 30.0, 20.0, 18.0, 15.0, 12.0, 10.0, 9.0, 7.0, 6.0, 5.0, 4.0, 3.0 };
float speedArrayInter[] = { 48.0, 30.0, 28.0, 24.0, 20.0, 18.0, 15.0, 12.0, 10.0, 9.0, 7.0, 6.0, 5.0, 4.0, 3.0, 2.4, 1.2, 0.24, 0.12 };

float sm90SpeedArrayIntra[] = { 60.0, 50.0, 40.0, 30.0, 24.0, 20.0, 15.0, 12.0, 11.0, 6.0, 3.0 };
float sm90SpeedArrayInter[] = { 48.0, 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 };

搜索算法使用贪心法和回溯法，整体而言一开始根据ncclTopoSystem中的maxBw和totalBw从速度数组里挑选一个最接近的数值作为搜索的开始，如果以此条件能搜索出结果则尝试再提升速度进行搜索，否则逐渐放松搜索要求进行搜索。可以发现代码中有大量的goto search语句，该语句前即是各类不同的搜索条件调整(搜索不到则降低条件，搜索到则提高条件)。

ncclResult_t ncclTopoCompute(ncclTopoSystem* system, struct ncclTopoGraph* graph) {
  int ngpus = system->nodes[GPU].count;
  
  // crossNic require more than 1 NET,
  // pattern need to be NCCL_TOPO_PATTERN_RING || NCCL_TOPO_PATTERN_BALANCED_TREE || NCCL_TOPO_PATTERN_SPLIT_TREE,
  // could be disabled using ncclParamCrossNic
  int crossNic = (system->nodes[NET].count > 1) &&
	 (graph->pattern == NCCL_TOPO_PATTERN_RING ||
          graph->pattern == NCCL_TOPO_PATTERN_BALANCED_TREE ||
          graph->pattern == NCCL_TOPO_PATTERN_SPLIT_TREE) ? ncclParamCrossNic() : 0;
  
  graph->crossNic = crossNic == 1 ? 1 : 0;
  graph->bwIntra = graph->bwInter = 0;
  graph->latencyInter = 0;
  
  // init with fastest intra connection type, will slow it down when searching
  graph->typeIntra = ngpus == 1 ? PATH_LOC : PATH_NVL; 
  
  // init with fastest inter connection type, will slow it down when searching
  graph->typeInter = PATH_PIX;
  
  graph->nChannels = 0;
  
  // Only try Heterogeneous channels in when connected wiht NVLS maybe because NVLS is configable
  int trySameChannels = graph->pattern == NCCL_TOPO_PATTERN_NVLS ? 0 : 1;
  graph->sameChannels = trySameChannels;

  int cpuArch, cpuVendor, cpuModel;
  NCCLCHECK(ncclTopoCpuType(system, &cpuArch, &cpuVendor, &cpuModel));

  // First try to load a graph from an XML file, search result can be cached in an XML file
  const char* str = ncclGetEnv("NCCL_GRAPH_FILE");
  if (str) {
    INFO(NCCL_ENV, "NCCL_GRAPH_FILE set by environment to %s", str);
    struct ncclXml* xml;
    NCCLCHECK(xmlAlloc(&xml, NCCL_GRAPH_XML_MAX_NODES));
    NCCLCHECK(ncclTopoGetXmlGraphFromFile(str, xml));
    int nChannels;
    NCCLCHECK(ncclTopoGetGraphFromXml(xml->nodes, system, graph, &nChannels));
    INFO(NCCL_GRAPH, "Search %d : %d channels loaded from XML graph", graph->id, nChannels);
    free(xml);
    if (graph->nChannels > 0) return ncclSuccess;
  }

  int ccMin;
  NCCLCHECK(ncclTopoGetCompCap(system, &ccMin, NULL));
  
  if (graph->pattern == NCCL_TOPO_PATTERN_NVLS && (system->nodes[NVS].count == 0 || ccMin < 90)) return ncclSuccess;
  
  // NVLS search must have ngpus heads at most.
  // Think about fully connected case, each GPU has nGPU links and each channel use 1 link of the GPU
  if (graph->pattern == NCCL_TOPO_PATTERN_NVLS) graph->maxChannels = system->nodes[GPU].count;

  // No intra nodes, use NCCL_TOPO_PATTERN_TREE
  if (ngpus == 1) if (graph->pattern != NCCL_TOPO_PATTERN_RING) graph->pattern = NCCL_TOPO_PATTERN_TREE;

  // force minChannels == maxChannels to force maximum channels searching to ensure maximum pallalism
  if (system->nodes[NET].count == 0 && graph->pattern == NCCL_TOPO_PATTERN_NVLS) {
    // Force intra-node NVLS algorithm to pull evenly from all GPUs.
    graph->minChannels = graph->maxChannels = system->nodes[GPU].count;
  }

  // save graph data for traceback algorithm to recovery later
  struct ncclTopoGraph tmpGraph;
  memcpy(&tmpGraph, graph, sizeof(struct ncclTopoGraph));

  // First try crossnic, then decrease bw and finally increase bwIntra.
  // Initialize bwIntra, bwInter using speedArray, choose the maximum speed close to system maxBw and totalBw.
  // Initialize gobalTimeout, each search will consume time based on different pattern.
  int nspeeds = 0;
  float* speedArray = NULL;
  if (system->nodes[NET].count == 0) {
    nspeeds = ccMin >= 90 ? NSPEEDSINTRA_SM90 : NSPEEDSINTRA;
    speedArray = ccMin >= 90 ? sm90SpeedArrayIntra : speedArrayIntra;
  } else {
    nspeeds = ccMin >= 90 ? NSPEEDSINTER_SM90 : NSPEEDSINTER;
    speedArray = ccMin >= 90 ? sm90SpeedArrayInter : speedArrayInter;
  }
  int pass = 1;
  int speedIndex = 0;
  float maxBw = system->maxBw;
  float totalBw = system->totalBw;
  if (ngpus > 1 && graph->pattern != NCCL_TOPO_PATTERN_RING) totalBw *= ngpus*1.0/(ngpus-1);
  while ((speedArray[speedIndex] > maxBw || speedArray[speedIndex]*graph->minChannels > totalBw) && speedIndex < nspeeds-1) speedIndex++;
  tmpGraph.bwIntra = tmpGraph.bwInter = speedArray[speedIndex];
  int64_t globalTimeout = NCCL_SEARCH_GLOBAL_TIMEOUT;

search:
  int time = tmpGraph.sameChannels ? NCCL_SEARCH_TIMEOUT_SAMECHANNELS :
    tmpGraph.pattern == NCCL_TOPO_PATTERN_TREE ? NCCL_SEARCH_TIMEOUT_TREE : NCCL_SEARCH_TIMEOUT;
  tmpGraph.nChannels = 0;
  globalTimeout -= time; 

  NCCLCHECK(ncclTopoSearchRec(system, &tmpGraph, graph, &time));
  
  // Optimal solution, stop here
  if (time == -1) goto done;
  if (graph->nChannels*graph->bwInter >= system->totalBw) goto done;

  if (pass == 1) {
    // First pass, we don't have a solution yet ; try other options
    // Try having different channels (except when going through AMD CPUs)
    if (tmpGraph.sameChannels == 1 &&
        !(cpuArch == NCCL_TOPO_CPU_ARCH_X86 && cpuVendor == NCCL_TOPO_CPU_VENDOR_AMD && tmpGraph.typeIntra == PATH_SYS)]{
      tmpGraph.sameChannels = 0;
      goto search;
    }
    tmpGraph.sameChannels = trySameChannels;

    if (time != -1) globalTimeout += time; // No result, traceback and search again
    else globalTimeout = NCCL_SEARCH_GLOBAL_TIMEOUT; // has
    if (globalTimeout < 0 && graph->nChannels) goto done;

    // Try a simpler tree
    if (ccMin >= 90 && tmpGraph.pattern == NCCL_TOPO_PATTERN_BALANCED_TREE) {
      tmpGraph.pattern = NCCL_TOPO_PATTERN_TREE;
      goto search;
    }
    tmpGraph.pattern = graph->pattern;

    int maxTypeIntra = system->nodes[NET].count > 0 ? tmpGraph.typeInter : PATH_SYS;
    if (tmpGraph.typeIntra < maxTypeIntra && (graph->nChannels == 0 || tmpGraph.typeIntra < graph->typeIntra)) {
      tmpGraph.typeIntra += 1;
      goto search;
    }
    tmpGraph.typeIntra = ngpus == 1 ? PATH_LOC : PATH_NVL;

    if (system->nodes[NET].count > 0 && tmpGraph.typeInter < PATH_SYS && (graph->nChannels == 0 || tmpGraph.typeInter < graph->typeInter || tmpGraph.typeInter < PATH_PXN)) {
      tmpGraph.typeInter += 1;
      goto search;
    }
    tmpGraph.typeInter = PATH_PIX;

    if (crossNic == 2 && tmpGraph.crossNic == 0) {
      // Try again with crossNic if permitted
      tmpGraph.crossNic = 1;
      goto search;
    }
    tmpGraph.crossNic = crossNic == 1 ? 1 : 0;

    // Decrease bw until we find a solution
    if ((speedIndex < nspeeds-1) && (graph->nChannels == 0 || (speedArray[speedIndex+1]/graph->bwInter > .49))) {
      tmpGraph.bwInter = tmpGraph.bwIntra = speedArray[++speedIndex];
      goto search;
    }
    speedIndex = 0;
    while (speedArray[speedIndex] > maxBw && speedIndex < nspeeds-1) speedIndex++;
    tmpGraph.bwIntra = tmpGraph.bwInter = speedArray[speedIndex];
  }

done:
  // We have a solution. Start from that solution and move to pass 2.
  if (pass == 1) {
    time = -1;
    NCCLCHECK(ncclTopoDupChannels(graph, ccMin, ngpus));
    memcpy(&tmpGraph, graph, sizeof(tmpGraph));
    speedIndex = 0;
    while (speedArray[speedIndex] > graph->bwInter && speedIndex < nspeeds-1) speedIndex++;
    tmpGraph.bwIntra = tmpGraph.bwInter = speedArray[speedIndex];
    tmpGraph.minChannels = graph->nChannels;
    pass = 2;
  }

  if (pass == 2) {
    // See if we can increase bw
    if (time != 0 && speedIndex > 0) {
      if (graph->pattern == NCCL_TOPO_PATTERN_RING) {
        // increase bw for Ring
        tmpGraph.bwIntra = tmpGraph.bwInter = speedArray[--speedIndex];
        goto search;
      } else if (graph->pattern == NCCL_TOPO_PATTERN_NVLS && tmpGraph.bwInter == graph->bwInter && tmpGraph.bwInter < tmpGraph.bwIntra*2) {
        tmpGraph.minChannels = tmpGraph.maxChannels = graph->nChannels;
        tmpGraph.bwInter = speedArray[--speedIndex];
        goto search;
      } else if (tmpGraph.bwIntra == graph->bwIntra && tmpGraph.bwIntra < tmpGraph.bwInter*2) {
        // increase bwIntra for trees (2 nodes or collnet)
        tmpGraph.bwIntra = speedArray[--speedIndex];
        goto search;
      }
    }
    time = -1;
    memcpy(&tmpGraph, graph, sizeof(tmpGraph));
  }

  if (graph->nChannels == 0 && graph->collNet == 0 && graph->pattern != NCCL_TOPO_PATTERN_NVLS) {
    WARN("Could not find a path for pattern %d, falling back to simple order", graph->pattern);
    for (int i=0; i<ngpus; i++) graph->intra[i] = system->nodes[GPU].nodes[i].gpu.rank;
    graph->inter[0] = graph->inter[1] = 0;
    graph->bwIntra = graph->bwInter = 0.1;
    graph->typeIntra = graph->typeInter = PATH_SYS;
    graph->nChannels = 1;
  }
  return ncclSuccess;
}

ncclTopoSearchRec

ncclTopoSearchRec执行ncclTopoCompute触发的每一次搜索。根据是否有网络节点，ncclTopoSearchRec内部通过调用ncclTopoSearchRecNet或者ncclTopoSearchTryGpu展开搜索。

/* Search Patterns
 *
 *     Intra-node
 * Ring            : GPU a -> GPU b -> .. -> GPU x -> GPU a
 * (=Split Tree Loop)
 * Tree            : GPU a -> GPU b -> .. -> GPU x
 * (=Split Tree)
 *
 *     Inter-node
 * Ring            : NET n -> GPU a -> GPU b -> .. -> GPU x -> NET n (or m if crossNic)
 * Tree            : NET n -> GPU a -> GPU b -> .. -> GPU x
 *                              `--> NET n (or m if crossNic)
 * Split Tree      : NET n -> GPU a -> GPU b -> .. -> GPU x
 *                                       `--> NET n (or m if crossNic)
 * Split Tree Loop : NET n -> GPU a -> GPU b -> .. -> GPU x -> GPU a
 *                                       `--> NET n (or m if crossNic)
 */
ncclResult_t ncclTopoSearchRec(struct ncclTopoSystem* system, struct ncclTopoGraph* graph, struct ncclTopoGraph* saveGraph, int* time) {
  // Initialize backToNet and backToFirstRank based on the search criteria
  int backToNet, backToFirstRank;
  NCCLCHECK(ncclTopoSearchParams(system, graph->pattern, &backToNet, &backToFirstRank));

  if (system->nodes[NET].count) {
    // Start from NET
    ncclTopoSearchRecNet(system, graph, saveGraph, backToNet, backToFirstRank, time);
  } else {
    // Intra-node only.
    if (graph->pattern == NCCL_TOPO_PATTERN_NVLS) {
      NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, 0, backToNet, backToFirstRank, 0, time, -1, -1, graph->nChannels));
      return ncclSuccess;
    } else if (graph->nChannels == 0) {
      // Try PCI order first
      NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, 0, backToNet, backToFirstRank, FORCED_ORDER_PCI, time, -1, -1, 0));
    } else {
      // Also try to replay previous channel
      int g;
      NCCLCHECK(ncclTopoReplayGetGpu(system, graph, -1, &g));
      NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, 0, backToNet, backToFirstRank, FORCED_ORDER_REPLAY, time, -1, -1, g));
    }
    if (graph->sameChannels == 0 || graph->nChannels == 0) {
      // Finally, try all other possibilities unless we are forced to use the same channels
      for (int g=0; g<system->nodes[GPU].count; g++) {
        NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, 0, backToNet, backToFirstRank, 0, time, -1, -1, g));
      }
    }
  }
  return ncclSuccess;
}

ncclTopoSearchRecNet

ncclTopoSearchRecNet是有NET节点时的搜索入口。该函数将graph->inter[graph->nChannels*2]设置好后启动ncclTopoSearchTryGpu。

ncclResult_t ncclTopoSearchRecNet(struct ncclTopoSystem* system, struct ncclTopoGraph* graph, struct ncclTopoGraph* saveGraph, int backToNet, int backToFirstRank, int* time) {
  const int bw = graph->bwInter;
  int* nets;
  NCCLCHECK(ncclCalloc(&nets, system->nodes[NET].count));
  int netCount;
  NCCLCHECK(ncclTopoSelectNets(system, graph->typeInter, -1, nets, &netCount));
  for (int i=0; i<netCount; i++) {
    if (graph->pattern == NCCL_TOPO_PATTERN_NVLS && i>0) continue; // skip other NETs for NVLS
    int n = nets[(graph->nChannels+i)%netCount];
    struct ncclTopoNode* net = system->nodes[NET].nodes+n;
    if (graph->collNet && net->net.collSupport == 0) continue; // skip NETs without collectives support if graph->collNet is set
    if (net->net.bw < bw) continue; // skip NETs with insufficient bandwidth
    if (graph->crossNic && (graph->nChannels & 1) && net->id != graph->inter[(graph->nChannels-1)*2+1]) continue; // if crossNic, in and out must be on different NICs

    graph->inter[graph->nChannels*2] = net->id;
    graph->latencyInter = net->net.latency;

    for (int i=0; i<system->nodes[NET].count; i++) {
      if ((system->nodes[NET].nodes[i].net.asic == net->net.asic) &&
          (system->nodes[NET].nodes[i].net.port == net->net.port)) {
        system->nodes[NET].nodes[i].net.bw -= bw;
      }
    }

    if (graph->pattern == NCCL_TOPO_PATTERN_NVLS) {
      // 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));
        if (gpu != -1) NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, 0, backToNet, backToFirstRank, 0, time, NET, n, gpu));
      }
    } else {
      if (graph->nChannels > 0) {
        // 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) {
        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
          int t = 1 << 10;
          NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, 0, backToNet, backToFirstRank, FORCED_ORDER_PCI, &t, NET, n, 0));
          if (t == -1) *time = -1;
        }

        // Then try the most local GPUs
        float maxBw = 0;
        int minHops = 0xfffffff;
        struct ncclTopoLinkList* paths = net->paths[GPU];
        for (int g=0; g<system->nodes[GPU].count; g++) {
          if (paths[g].bw > maxBw) {
            maxBw = paths[g].bw;
            minHops = paths[g].count;
          } else if (paths[g].bw == maxBw && paths[g].count < minHops) {
            minHops = paths[g].count;
          }
        }
        if (maxBw >= bw) {
          for (int i=0; i<system->nodes[GPU].count; i++) {
            int g = (graph->nChannels+i)%system->nodes[GPU].count;
            if (paths[g].bw == maxBw && paths[g].count == minHops) {
              NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, 0, backToNet, backToFirstRank, 0, time, NET, n, g));
            }
          }
        }
      }
    }

    for (int i=0; i<system->nodes[NET].count; i++) {
      if ((system->nodes[NET].nodes[i].net.asic == net->net.asic) &&
          (system->nodes[NET].nodes[i].net.port == net->net.port)) {
        system->nodes[NET].nodes[i].net.bw += bw;
      }
    }
  }
  free(nets);
  return ncclSuccess;
}

ncclTopoSearchTryGpu

ncclTopoSearchTryGpu是搜索GPU节点时的搜索入口。该函数调用递归函数ncclTopoSearchRecGpu完成搜索过程。

ncclResult_t ncclTopoSearchTryGpu(struct ncclTopoSystem* system, struct ncclTopoGraph* graph, struct ncclTopoGraph* saveGraph, int step, int backToNet, int backToFirstRank, int forcedOrder, int *time, int type, int index, int g) {
  const uint64_t flag = 1ULL<<(graph->nChannels);
  struct ncclTopoNode* gpu;
  NCCLCHECK(ncclTopoFollowPath(system, graph, type, index, GPU, g, 1, &gpu));
  if (gpu) {
    gpu->used ^= flag;
    NCCLCHECK(ncclTopoSearchRecGpu(system, graph, saveGraph, gpu, step, backToNet, backToFirstRank, forcedOrder, time));
    gpu->used ^= flag;
    NCCLCHECK(ncclTopoFollowPath(system, graph, type, index, GPU, g, -1, &gpu));
  }
  return ncclSuccess;
}

ncclTopoSearchRecGpu

ncclTopoSearchTryGpu逻辑的主要实现在ncclTopoSearchRecGpu中。该函数不断递归搜索，递归终止条件是step==ngpu，即需要遍历所有GPU节点。当终止条件到达时，通过ncclTopoCompareGraphs函数判断是否找到更优的方案，如果graph->nChannels < graph->maxChannels此时针对新通道展开新一轮搜索。搜索过程中，不同的step处理方式不同，一般而言便是不断寻找下一个GPU，但比如遇到step == backToFirstRank等特殊情况则需要找到初始的GPU节点完成环的回路。在搜索推进过程中会不断调用ncclTopoFollowPath函数进行剪枝，该函数判断下一跳是否会引入带宽瓶颈，如果不会则继续往前搜索，否则可以提前结束搜索。

ncclResult_t ncclTopoSearchRecGpu(struct ncclTopoSystem* system, struct ncclTopoGraph* graph, struct ncclTopoGraph* saveGraph, struct ncclTopoNode* gpu, int step, int backToNet, int backToFirstRank, int forcedOrder, int *time) {
  if ((*time) <= 0) return ncclSuccess;
  (*time)--;

  int ngpus = system->nodes[GPU].count;
  if (step == ngpus) {
    // Determine whether we found a better solution or not
    int copy = 0;
    graph->nChannels++;
    NCCLCHECK(ncclTopoCompareGraphs(system, graph, saveGraph, &copy));
    if (copy) {
      memcpy(saveGraph, graph, sizeof(struct ncclTopoGraph));
      if (graph->nChannels == graph->maxChannels) *time = -1;
    }
    if (graph->nChannels < graph->maxChannels) {
      // not reaching maxChannels, continue to search
      NCCLCHECK(ncclTopoSearchRec(system, graph, saveGraph, time));
    }
    graph->nChannels--;
    return ncclSuccess;
  }
  graph->intra[graph->nChannels*ngpus+step] = gpu->gpu.rank;
  int g = gpu - system->nodes[GPU].nodes;
  if (step == backToNet) {
    // first get back to NIC
    if (system->nodes[NET].count) {
      int startNetIndex;
      NCCLCHECK(getNetIndex(system, graph->inter[graph->nChannels*2], &startNetIndex));
      struct ncclTopoNode* startNet = system->nodes[NET].nodes+startNetIndex;
      int netCount;
      int* nets;
      NCCLCHECK(ncclCalloc(&nets, system->nodes[NET].count));
      NCCLCHECK(ncclTopoSelectNets(system, graph->typeInter, g, nets, &netCount));
      for (int i=0; i<netCount; i++) {
        int n = nets[i];
        struct ncclTopoNode* net = system->nodes[NET].nodes+n;
        if (graph->pattern == NCCL_TOPO_PATTERN_TREE && net->id != startNet->id) continue; // Trees are symmetric
        if (graph->crossNic != 1 && (net->net.asic != startNet->net.asic || net->net.port != startNet->net.port)) continue;
        if (graph->crossNic && (graph->nChannels & 1) && net->id != graph->inter[(graph->nChannels-1)*2]) continue;

        // Balanced Tree : count half of the bandwidth on first two GPUs
        int nextBackToNet = -1;
        float bwInterSave = graph->bwInter;
        if (graph->pattern == NCCL_TOPO_PATTERN_BALANCED_TREE) {
          // Count half of the bandwidth on each of the first two GPUs
          if (step == 0) nextBackToNet = 1;
          else if (net->id != graph->inter[graph->nChannels*2+1]) continue;
          graph->bwInter /= 2;
        }

        NCCLCHECK(ncclTopoFollowPath(system, graph, GPU, g, NET, n, 1, &net));
        graph->bwInter = bwInterSave;
        if (net) {
          graph->inter[graph->nChannels*2+1] = net->id;
          NCCLCHECK(ncclTopoSearchRecGpu(system, graph, saveGraph, gpu, step, nextBackToNet, backToFirstRank, forcedOrder, time));

          if (graph->pattern == NCCL_TOPO_PATTERN_BALANCED_TREE) graph->bwInter /= 2;
          NCCLCHECK(ncclTopoFollowPath(system, graph, GPU, g, NET, n, -1, &net));
          graph->bwInter = bwInterSave;
        }
      }
      free(nets);
    }
  } else if (graph->pattern == NCCL_TOPO_PATTERN_NVLS) {
    NCCLCHECK(ncclTopoSearchTryNvls(system, graph, saveGraph, g, ngpus, time));
  } else if (step < system->nodes[GPU].count-1) {
    // Go to next GPU
    int next[NCCL_TOPO_MAX_NODES];
    int count;
    if (forcedOrder == FORCED_ORDER_PCI) { // Try the PCI order
      next[0] = step+1;
      count = 1;
    } else if (forcedOrder == FORCED_ORDER_REPLAY) { // Try last channel order
      NCCLCHECK(ncclTopoReplayGetGpu(system, graph, step, next));
      count = 1;
    } else { // Normal search
      NCCLCHECK(ncclTopoSearchNextGpuSort(system, graph, gpu, next, &count, backToNet == -1 ? 0 : backToNet == step+1 ? 1 : -1 ));
    }
    for (int i=0; i<count; i++) {
      NCCLCHECK(ncclTopoSearchTryGpu(system, graph, saveGraph, step+1, backToNet, backToFirstRank, forcedOrder, time, GPU, g, next[i]));
    }
  } else if (step == backToFirstRank) {
    // Find first GPU and loop back to it
    int p;
    NCCLCHECK(getGpuIndex(system, graph->intra[graph->nChannels*ngpus], &p));
    struct ncclTopoNode* firstGpu;
    NCCLCHECK(ncclTopoFollowPath(system, graph, GPU, g, GPU, p, 1, &firstGpu));
    if (firstGpu) {
      NCCLCHECK(ncclTopoSearchRecGpu(system, graph, saveGraph, firstGpu, step+1, backToNet, -1, forcedOrder, time));
      NCCLCHECK(ncclTopoFollowPath(system, graph, GPU, g, GPU, p, -1, &firstGpu));
    }
  } else {
    // Next path
    NCCLCHECK(ncclTopoSearchRecGpu(system, graph, saveGraph, gpu, ngpus, -1, -1, forcedOrder, time));
  }
  return ncclSuccess;
}

ncclTopoFollowPath

ncclTopoFollowPath是搜索过程中的剪枝函数，其寻找type1/index1节点至type2/index2节点的路径类型和带宽是否满足ncclTopoGraph中typeintra，bwIntra，bwInter以及typeinter的要求。首先排除路径类型不匹配的情况（比如graph要求NVL而路径却是PCI），如路径类型判断通过，还需调用followPath判断路径带宽是否满足要求。

// Try to go from node type1/index1 to no type2/index2. mult indicates whether we are counting the bandwidth (1) or undoing (-1).
// Set node to the destination node if successful, NULL otherwise.
static ncclResult_t ncclTopoFollowPath(struct ncclTopoSystem* system, struct ncclTopoGraph* graph, int type1, int index1, int type2, int index2, int mult, struct ncclTopoNode** node) {
  // First handle easy cases
  *node = system->nodes[type2].nodes+index2;
  if (type1 == -1) return ncclSuccess;
  struct ncclTopoNode* node1 = system->nodes[type1].nodes+index1;
  struct ncclTopoLinkList* path = node1->paths[type2]+index2;
  struct ncclTopoNode* node2 = system->nodes[type2].nodes+index2;
  struct ncclTopoLinkList* revPath = node2->paths[type1]+index1;

  if (path == NULL) {
    WARN("No path computed to go from %s/%d to %s/%d", topoNodeTypeStr[type1], index1, topoNodeTypeStr[type2], index2);
    return ncclInternalError;
  }

  // Now check link type
  *node = NULL;
  int intra = (type1 == GPU || type1 == NVS) && (type2 == GPU || type2 == NVS);
  float bw = intra ? graph->bwIntra : graph->bwInter;
  int type = intra ? graph->typeIntra : graph->typeInter;

  if (mult == 1 && (path->type > type)) return ncclSuccess;
  if (mult == 1 && (graph->pattern == NCCL_TOPO_PATTERN_BALANCED_TREE ||
        graph->pattern == NCCL_TOPO_PATTERN_TREE ||
        graph->pattern == NCCL_TOPO_PATTERN_SPLIT_TREE) &&
      (revPath->type > type)) return ncclSuccess;

  bw *= mult;

  // Check there is enough bandwidth on paths.
  int step = 0;
  NCCLCHECK(followPath(path, node1, path->count, bw, &step));
  if (step < path->count) goto rewind;

  // Enough bandwidth : return destination node.
  graph->nHops += mult*path->count;
  *node = system->nodes[type2].nodes+index2;
  return ncclSuccess;

rewind:
  // Not enough bandwidth : rewind and exit.
  NCCLCHECK(followPath(path, node1, step, -bw, &step));
  return ncclSuccess;
}

followPath

followPath函数便是判断一条路径带宽是否满足要求的函数。

// Check if there is enough bandwidth on the path starting from start node with maxSteps steps and bw bandwidth.
static ncclResult_t followPath(struct ncclTopoLinkList* path, struct ncclTopoNode* start, int maxSteps, float bw, int* steps) {
  float pciBw = bw;

  // Check if there is INTEL CPU in the path, if so, we need to consider the overhead of P2P between INTEL CPU
  for (int step=0; step<path->count; step++) {
    struct ncclTopoNode* node = path->list[step]->remNode;
    if (node->type == CPU) {
      // Account for P2P inefficiency through Intel CPU RC
      if (path->type == PATH_PHB && start->type == GPU &&
          node->cpu.arch == NCCL_TOPO_CPU_ARCH_X86 &&
          node->cpu.vendor == NCCL_TOPO_CPU_VENDOR_INTEL) {
        pciBw = INTEL_P2P_OVERHEAD(bw);
      }
    }
  }

  // Check if there is enough bandwidth on the path
  struct ncclTopoNode* node = start;
  for (int step=0; step<maxSteps; step++) {
    struct ncclTopoLink* link = path->list[step];
    struct ncclTopoLink* revLink = NULL;
    float fwBw = link->type == LINK_PCI ? pciBw : bw;

    // Update reverse bandwidth if needed, based on the type of the link and the type of the nodes
    float revBw = 0;
    if (link->remNode->type == GPU && link->remNode->gpu.cudaCompCap < 80 && start->type != GPU) {
      if (revLink == NULL) NCCLCHECK(findRevLink(node, link->remNode, link->type, &revLink));
      revBw += fwBw/8;
    }
    if (link->remNode->type == CPU && link->remNode->cpu.arch == NCCL_TOPO_CPU_ARCH_POWER && link->type == LINK_NVL) {
      if (revLink == NULL) NCCLCHECK(findRevLink(node, link->remNode, link->type, &revLink));
      revBw += fwBw;
    }
    // Check if there is enough forward bandwidth on the link and backward bandwidth on the reverse link, return the step if not
    if (link->bw < fwBw || (revBw && revLink->bw < revBw)) { *steps = step; return ncclSuccess; }
    SUB_ROUND(link->bw, fwBw);
    if (revBw) SUB_ROUND(revLink->bw, revBw);
    node = link->remNode;
  }

  // we have enough bandwidth on the whole path if we checked all steps
  *steps = maxSteps;
  return ncclSuccess;
}
慢河
关注
专栏目录
nccl-repo-ubuntu1804-2.4.2-ga-cuda10.1_1-1_amd64.deb
06-09
nccl
nccl-local-repo-ubuntu2004-2.9.9-cuda11.0_1.0-1_amd64.deb
06-23
nccl本地安装包，适合ubuntu20.04及cuca11.0以上版本
NCCL拓扑管理 - Topo模块
漫河
07-08 2153
Topo模块提供系统拓扑的构建接口，依赖XML模块构建系统拓扑图，根据实际情况合并Peer的系统拓扑信息用以构造超级节点（GPU南向互联），增加节点之间的连接，为后续Path模块的路径初始化做准备。
理解NCCL源码必看：详细论述NCCL源码做了哪些事
这个地瓜真是甜啊
03-21 3940
详解NCCL 源码的实现的功能以及如何实现的
理解NCCL的Tree
03-25 1888
结合了github中issue的解答，自己在环境上通过log一步一步把Tree里的关系画出来，发现自己好像明白了一些东西，整理了一下对Tree的理解
NCCL拓扑管理 - Path模块
漫河
07-10 346
NCCL Path模块提供接口为Topo模块产生的的系统拓扑图的节点提供路径初始化。Path模块使用贪心算法（构造最少跳数且最大带宽的路径）对系统拓扑图中每一个节点到其他节点的路径进行初始化，为后续Search模块的逻辑拓扑搜索做准备。
NVIDIA NCCL 源码学习（十四）- NVLink SHARP
KIDGIN7439的专栏
04-18 3815
上节我们介绍了IB SHARP的工作原理，进一步的，英伟达在Hopper架构机器中引入了第三代NVSwitch，就像机间IB SHARP一样，机内可以通过NVSwitch执行NVLink SHARP，简称nvls，这节我们会介绍下NVLink SHARP如何工作的。
【分布式】小白看Ring算法 - 03
参谋谋的博客
11-23 3085
NCCL（NVIDIA Collective Communications Library）是由NVIDIA开发的一种用于多GPU间通信的库。NCCL的RING算法是NCCL库中的一种通信算法，用于在多个GPU之间进行环形通信。RING算法的基本思想是将多个GPU连接成一个环形结构，每个GPU与相邻的两个GPU进行通信。数据沿着环形结构传递，直到到达发送方的位置。这样的环形结构可以有效地利用GPU之间的带宽，提高通信的效率。RING算法的步骤如下：fill:#333;color:#333;
nccl-local-repo-ubuntu2004-2.8.4-cuda11.1-1.0-1-amd64.zip
09-11
标题中的"nccl-local-repo-ubuntu2004-2.8.4-cuda11.1-1.0-1-amd64.zip"指的是NVIDIA Collective Communication Library (NCCL) 的本地仓库包，专为Ubuntu 20.04操作系统设计，并且与CUDA 11.1版本兼容。NCCL是一个...
nccl_2.11.4-1+cuda11.0_x86_64.txz
03-09
nccl_2.11.4-1+cuda11.0_x86_64.txz
nccl-repo-ubuntu1604-2.4.8-ga-cuda9.0_1-1_amd64.deb 并行计算nccl
12-17
tensorflow并行计算库文件nccl,nccl-repo-ubuntu1604-2.4.8-ga-cuda9.0_1-1_amd64.deb 适用于16.04ubuntu系统
使用c#制作一个小型桌面程序
qq_44691564的博客
09-14 3272
创建.NET Framework新项目，将之前生成的dll放在Demo文件夹的bin ->debug或是 release中（看你自己用的什么模式），创建完成后 系统会自动生成一些文件，其中 pch.cpp 先不要修改，pch.h中先导入自己需要用到的库，下面是我的代码。布局完了之后会自动生成Form1.Designer.cs 的窗口设计代码，点击控件按F4 还可以修改他们的属性。完成之后点击生成 就可以在bin中出现的你的.exe文件咯，是不是很简单呀~[狗头]我们还需要一个入口主程序。
零售业的数字化转型与消费者体验升级
最新发布
weixin_58820787的博客
09-24 703
在数字化浪潮的推动下，零售业正经历着前所未有的变革。数字化转型不仅为零售商带来了新的商业模式和运营效率的提升，更重要的是，它极大地提升了消费者的购物体验。金智维将探讨零售业如何通过数字化转型，实现线上线下融合、智能推荐、个性化服务等方面的升级，以满足消费者日益增长的需求。
介绍一下常用的激活函数？
GamBleout的博客
09-19 1234
死神经元：当神经网络中的某些神经元在训练过程中始终不被激活，即其输出值始终接近于零，这些神经元就被称为“死神经元”。与Sigmoid函数类似，也存在梯度消失问题，但相对于Sigmoid函数，其输出范围更大，有利于梯度的传播。Leaky ReLU是ReLU的一个变体，在输入为负时给予一个小的非零斜率，避免了“死神经元”问题。在正区间内解决了梯度消失问题，但在负区间内梯度为0，可能导致“死神经元”问题。将一组数值转换为概率分布，即所有输出的和为1，且每个输出都在(0,1)之间。当输入为负时，输出为0。
苹果AI手机遇阻，国产手机找到超车机遇
songguocaijing的博客
09-23 901
去年，它只提供12期免息分期。事实上，苹果的Apple Intelligence高度依赖于Open AI，然而Open AI在国内的使用受限，而目前苹果在中国的AI合作方一直尚未决定，此前有消息报道，苹果已与多家中国大模型公司进行了接触，包括百度、阿里巴巴以及百川智能等，但也没有明确消息流出。苹果的光环总是高居天上，今年6月10日，苹果发布全新智能体Apple Intelligence后，市场对于将搭载“Apple所开发出的独一无二的AI”新机型寄予了极高的期待，甚至激起了对苹果下一个“超级周期”的预期。
IT行业中的技术趋势与未来展望
2301_78944833的博客
09-21 2002
IT行业正在经历一场前所未有的技术变革，各种新兴技术相互融合，推动着企业和社会的数字化转型。从云计算到人工智能，从物联网到区块链，这些技术趋势不仅改变了IT行业的面貌，也塑造了未来世界的形态。无论是IT从业者还是企业管理者，紧跟这些技术趋势，将在未来的竞争中占据有利位置。
AI下凡记：智能体如何走进我们的日常生活
Enweitech Software Works
09-19 1177
随着AI行业的不断发展，未来必然会出现更自然的人机交互方式，智能体将会充当中间的沟通桥梁，App将解构成颗粒度更小的服务，由智能体担当智能耦合的角色，在这样的背景下，智能体将有望成为新的流量入口，所以，提前布局智能体，对大厂的商业生态非常重要。因此，新一轮的智能体争夺战，虽然从流量入口开始，但最终考验的却是企业的生态运营能力，企业前端要有足够的场景化牵引力，内部要有流畅的场景协调能力，外部要有创新的模式，以及提供持续优化的用户体验，才是智能体时代优秀的“流量入口”。后者则是创造需求，引来流量。
(chat) ubuntu@ubun:~$ sudo apt-get update 获取:1 file:/var/nccl-local-repo-ubuntu2004-2.14.3-cuda11.7 InRelease [1,575 B] 获取:1 file:/var/nccl-local-repo-ubuntu2004-2.14.3-cuda11.7 InRelease [1,575 B] 错误:1 file:/var/nccl-local-repo-ubuntu2004-2.14.3-cuda11.7 InRelease 由于没有公钥，无法验证下列签名： NO_PUBKEY 4FA3634244000BE4 命中:2 http://ftp.sjtu.edu.cn/ubuntu focal InRelease 命中:3 http://ftp.sjtu.edu.cn/ubuntu focal-security InRelease 命中:4 http://security.ubuntu.com/ubuntu focal-security InRelease 命中:5 https://developer.download.nvidia.cn/compute/cuda/repos/ubuntu2004/x86_64 InRelease 命中:6 http://ppa.launchpad.net/graphics-drivers/ppa/ubuntu focal InRelease 正在读取软件包列表... 完成 W: GPG 错误：file:/var/nccl-local-repo-ubuntu2004-2.14.3-cuda11.7 InRelease: 由于没有公钥，无法验证下列签名： NO_PUBKEY 4FA3634244000BE4 E: 仓库 “file:/var/nccl-local-repo-ubuntu2004-2.14.3-cuda11.7 InRelease” 没有数字签名。 N: 无法安全地用该源进行更新，所以默认禁用该源。 N: 参见 apt-secure(8) 手册以了解仓库创建和用户配置方面的细节。
06-13
这个错误提示是因为在更新软件包列表时，apt-get无法验证`/var/nccl-local-repo-ubuntu2004-2.14.3-cuda11.7`这个仓库的数字签名。你可以按照以下步骤来解决这个问题： 1. 打开终端，输入以下命令来获取缺失的公钥： ``` sudo apt-key adv --keyserver keyserver.ubuntu.com --recv-keys 4FA3634244000BE4 ``` 其中，`4FA3634244000BE4`是缺失的公钥的ID。 2. 输入以下命令来更新软件包列表： ``` sudo apt-get update ``` 如果一切正常，这时你应该可以成功更新软件包列表了。 希望这可以帮助到你。
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值