NVIDIA GPGPU通信架构

admin管理员组

文章数量:1607884

NVIDIA GPGPU通信架构

吃果冻不吐果冻皮 2024年07月18日 11:50 四川

以下文章来源于大模型新视界 ，作者Bruce 仗剑走天涯

大模型新视界.

大模型产品、算法、应用

写在前面

此一部分，介绍nvlink、nccl、nvswitch、gpgpu 之间的具体联动关系。侧重于介绍通信系统本身，以及和计算的耦合性，Nvlink & nvswitch system 为 nvidia gpgpu 大规模计算和超大算力的支持可以说是十分重要的，为此有必要单列一章进行分析介绍。

NVLink 之前

传统互联通路的失灵 - PCIE太慢了

PCIe 技术是PCI 技术的扩展，最初由Intel 于2001年公布，原名3GIO（第三代IO），在2002年由PCI-SIG 审核通过，改名为PCI express，此后，每一代带宽都是前一代的2倍，PCIE gen 5 每lane为32Gbps，折合为3938MB/s，gen 5 x16 即为64GB/s。2022年公布PCie gen 6 specification，每lane 带宽为64Gbps，x16 可达1Tbps 以上。

第一代Nvlink 起于pascal，2016年，Pcie 3.0 时期，彼时，gen3 x16的带宽只有128Gbps，也就是16GB/s，而且是机器上所有的设备只能共享这个带宽，而显然对于一机8卡的GPU 高性能计算需求，是远远不够的，如果按GPU 独占带宽，且8卡均衡带宽计算，每卡不过16Gbps（2GB/s），何况NIC 网卡本身也是一个PCIE 的带宽的消费者，算上带宽竞争，则更加寥寥。

P 100 提出的带宽需求是多少？是什么促使NVIDIA 彻底抛弃PCIE，走向了自研高速互联协议呢？

我们查看白皮书，P100的32位浮点计算能力为10.6TFLOPS，也就是说每秒执行10.6 * 10^12 次32位浮点计算，如果每个浮点都是GPU外搬运而来且完全不复用旧数据，则需要10.6 * 10 ^ 3 * 32 Gbps的带宽，40000 GB/s，当然实际上不会如此，数据会有一定程度的复用，这和业务相关。不过显然每卡2GB/s的带宽能力，相差20000倍，约等于用户基于1B 及其结果原地计算20000次，要求用户达到这么高的数据复用是不太现实的。
我们比较一下，2016年推出的nvlink v1 实现了多大的带宽呢？
每GPU 有4个nvlink，每个link 单向20GB/s，每卡双向总带宽达160GB/s，这就显然比Pcie 3 x16 提升了2个数量级的带宽，对用户数据复用的要求也大大降低了。

RDMA 网络的介入 - IB/Mellanox products家族

除去传统的PCIE 通路外，CPU 和 传统的TCP 通路也逐渐显现出GPU 训练的带宽窘境，这比Pcie 带宽出现的更早，在2012年，NVIDIA 在kepler 上就推出了GPUdirect的能力，希望数据传输上bypass CPU，且特别指出了RDMA 的使用是对于GPUDirect 的高性能的保证，后来该能力更多被直接称为GPUDirectRDMA，简称GDR。GDR 至今已经10年了，10年前GPU 高性能领域就已经开始慢慢淘汰TCP，而逐步使用RDMA NIC 作为高性能数据计算的基础网卡。

至2020年，Amphere 支持Mellanox NIC 和 相关产品，NVIDIA 收购Mellanox，GDR带宽都在稳步发展，rdma 网络带宽能力，逐步从10Gbps 发展到如今CX 7的400Gbps。
我们也可以简单看一下IB 规格的历史。

虽然2012年的rdma 还是10Gbps的时代，但已经占据了超级计算机互联的绝大部分，至2009年，世界500强超级计算机中，259个使用千兆以太网作为内部互连技术，181个使用InfiniBand。同期的TCP 带宽暂未找到相关数据。
随着cable 性能的提高，rdma/ib 带宽一路高歌猛进，至今CX7 为NDR * 4 = 400Gbps。而2016年，mellanox 率先提出了在网计算协议SHARP，在自研交换机上率先实现了aggregation reduction，也就是交换机实现reduce 计算，简化网络包数的技术。至2020年，Nvidia 收购 mellanox，sharp v2 已经在mellanox 交换机上实现，至今已实现sharpv3且在quantum 中使用，而NCCL 开始支持sharp 的时间是2019年。具体sharp 的原理与收益我们会在之后的章节说明。

NV-SLI /SLI bridge

在提nvlink 之前，我们还需要了解SLI。一种早期的scaleble interface，SLI 用于早期多显卡互联，SLI 互联的GPU，可以共享存储和计算任务，SLI 互联的桥接部件称为SLI Bridge，简图如下。也是从SLI 为起点，诞生出的nvlink，后来的nvlink 可以认为是SLI的升级版本。至今，NVLINK 在Turing 也重新支持了SLI，但是性能相对低（支持路数较少），之后的nvlink 也不再提到SLI的场景，因此怀疑可能和向前兼容有关。

高速互联设备 - NVLink & NVSwitch

NVLINK

Nvlink 自从pascal 架构提出以来，至今已到第四代。每lane的带宽以及每link的带宽数都在增加。

V 1

主要特性

NVLink 1.0是为GPU-GPU、GPU-CPU高速互连的接口，支持直接读写对端CPU/GPU的内存（所有内存都在共享地址空间里）。主要特性：
1.每个link双向接口，每个方向由8 lane组成，单lane最高速率20Gbps，单link 单向带宽为20Gbps x8 = 20GBps，双向带宽40GBps。
2.单GPU（P100）支持4NVLink，双向带宽一共160GBps
3.提供load/store语义，让用户能对peer GPU内存进行read/writes操作，另外还支持atomics操作
4.NVLink 1.0是一种基于包的协议，包长在一定范围可变
5.不支持多队列，仅支持多VC（virtual channel）
6.Flow control：在请求包里带flow control credit
7.通过CRC检测数据错误
8.Replay: 类似Go-back-N重传
9.仅支持和部分CPU（IBM Power系列）互连。

协议

类似PCIe， NVLink 1.0也分为Physical Layer、Data Link Layer和Transaction Layer。

1.Physical Layer: 接PHY，负责deskew、framing、(de)scrambling、polarity inversion和lane reversal等
2.Data Link Layer: 负责可靠性传输，通过CRC/ACK等实现
3.Transaction Layer: 负责synchronization, link flow control, virtual channels， 并能将多个nvlink汇聚到一起

包格式

NVLink 1.0的数据包格式如图：

1.每次transaction至少包含一个request和一个response（Posted operation除外）
2.以128bit为一个基本单位（flit）,单个Packet支持1-18个flit，单个packet可能包含0-16个data flit， 即最多传输256B的data。
3.包头包含三个部分：

• CRC：校验前前包的header和上一个包的payload

• Header：包含request type, address, flow control credits和tag identifier

• DL Header：包含acknowledge identifier, packet length information和application number tag

1.AE（可选）：传输Comand-specific信息，或者用于修改Command的默认值，只有在变化时传输
2.BE（可选）：write和atomic指定需要写的字节，类似掩码

CRC和重传机制

1.如果CRC校验成功，则回positive ack，校验失败不回复ack
2.请求侧数据缓存在replay buffer
3.如果请求侧收到正确的ack sequence，则将packet从replay buffer删除
4.如果请求侧遇到错误的ack Sequence，或者遇到timeout，则回退到上一个Acked的包，从relpay buffer进行重传，即go back N。

和其它模块的接口

NVLink通过High Speed Hub和其它模块互连，HSHub连接到GPU的Crossbar、High Speed Copy Engine和PCIe等其它接口。可以看到copyengine 可以选择pcie 或者nvlink 进行copy功能。

吃果冻不吐果冻皮

专注于AI工程化（LLM、MLOps、LLMOps、RAG、Agent）落地。

148篇原创内容

公众号

基本拓扑构型

V2

与基于Volta推出，主要载体是V100，与1.0版本相比，主要区别在于：

1.NVLink2.0有更快的带宽，单NVLink带宽由40GBps提升至50GBps，单GPU有6个NVLink，一共300GBps
2.支持low-power operation mode
3.CPU 方向相关：Cache 一致性增强，支持CPU 通过nvlink 读取数据到cache，并支持了更加完善的GPU-CPU atomic；支持 ATS

注意，由于和IBM 的合作企图推出更强大的超级计算机，这一代对于CPU侧的增强较多，对于GPU 侧主要是堆规格。但另一方面，这一代出现了nvswitch 1.0。

V3

基于 Ampere 提出。 NVLink 3.0 的特点:
lossless, highbandwidth, low-latency shared memory interconnect, 并通过link level error detection、packet replay保证传输可靠性。NVLink 3.0的新特性包括：

1. 更高的带宽、信号线数减半、单GPU 12个nvlink

2. 升级error detection和recovery。

3. Write操作变成non-posted，使得请求侧可以进行同步，错误处理也有改进

4. 优化了small payload write和没有data的response的效率

V4

基于hopper，NVLink 4.0特性：

1.单个nvlink只用了2个lane实现单向25GBps。单个GPU支持18个nvlink，总共900GBps带宽，是上一代的1.5倍。
2.为了支持跨多个node的集群，引入了NVLink Network。

• 不再是所有GPU共享地址空间，引入Network Address Space，和GPU地址空间隔离，H100支持地址间的翻译

• NVLink 4.0和IB一样，user software需要先建立连接。

NVSwitch

Nvswitch 自从volta 架构提出以来，至今已到第三代。

V1

nvswitch是一块独立的nvlink芯片，其提供了高达18路nvlink的接口其支持nvlink 2.0，也就意味着每个接口均能提供双信道高达50GB/s的带宽，那么这块芯片总计能够提供900GB/s的带宽。这块芯片功率100w，基于台积电12nm FinFet FFN nvidia订制工艺，来源于增强的16nm节点，拥有2b个晶体管。

这块die封装在1940个pin大小为4cm2的BGA芯片中，其中576个针脚专门服务于18路的nvlink，剩下的阵脚则用于电源，或者其他I/O接口，比如用于管理端口的x4 pcie，I2c，GPIO等等。

通过nvswitch提供的18路接口，nvswitch能够让nvidia设计出完全无阻塞的全互联16路GPU系统。每块v100中的6路nvlink将分别连接到6块nvswitch上面。这样8块v100与6块nvsiwtch完全连接，构成一个基板。

V2

每个switch 18link。

V3

NVSwitch 3.0特性：

1.单芯片，64 port,12.8Tbps
2.可以在node内和node外放置
3.实现了collective operation的硬件加速，包括multicast和SHARP，加速了了allreduce，reduce_scatter
4.2跳交换，L1 Nvswitch 与 GPU 相连接口为nvlink，与L2 交换机连接为osfp。
5.每个nvswitch 芯片包含64个nvlink，每四个nvlink 组成一个osfp，故有16个osfp

NCCL 软件栈浅说

NCCL 本身很复杂，我这里推荐oneflow的NCCL 专栏。
https://mp.weixin.qq/mp/appmsgalbum

oneflow 的解析写得十分详细（抱拳），我这里仅做一点自己读源码的总结。

我们这里不讲一般NIC network相关的软件栈。主要集中于GPU和nvlink的链路。

基本例子

简单的集合通信例子

int main(int argc, char* argv[])
{
  int size = 32*1024*1024;
  int myRank, nRanks, localRank = 0;
  //initializing MPI
  MPICHECK(MPI_Init(&argc, &argv));
  MPICHECK(MPI_Comm_rank(MPI_COMM_WORLD, &myRank));
  MPICHECK(MPI_Comm_size(MPI_COMM_WORLD, &nRanks));
  //calculating localRank which is used in selecting a GPU
  uint64_t hostHashs[nRanks];
  char hostname[1024];
  getHostName(hostname, 1024);
  hostHashs[myRank] = getHostHash(hostname);
  MPICHECK(MPI_Allgather(MPI_IN_PLACE, 0, MPI_DATATYPE_NULL, hostHashs, sizeof(uint64_t), MPI_BYTE, MPI_COMM_WORLD));
  for (int p=0; p<nRanks; p++) {
     if (p == myRank) break;
     if (hostHashs[p] == hostHashs[myRank]) localRank++;
  }
  //each process is using two GPUs
  int nDev = 2;
  float** sendbuff = (float**)malloc(nDev * sizeof(float*));
  float** recvbuff = (float**)malloc(nDev * sizeof(float*));
  cudaStream_t* s = (cudaStream_t*)malloc(sizeof(cudaStream_t)*nDev);
  //picking GPUs based on localRank
  for (int i = 0; i < nDev; ++i) {
    CUDACHECK(cudaSetDevice(localRank*nDev + i));
    CUDACHECK(cudaMalloc(sendbuff + i, size * sizeof(float)));
    CUDACHECK(cudaMalloc(recvbuff + i, size * sizeof(float)));
    CUDACHECK(cudaMemset(sendbuff[i], 1, size * sizeof(float)));
    CUDACHECK(cudaMemset(recvbuff[i], 0, size * sizeof(float)));
    CUDACHECK(cudaStreamCreate(s+i));
  }
  ncclUniqueId id;
  ncclComm_t comms[nDev];
  //generating NCCL unique ID at one process and broadcasting it to all
  if (myRank == 0) ncclGetUniqueId(&id);
  MPICHECK(MPI_Bcast((void *)&id, sizeof(id), MPI_BYTE, 0, MPI_COMM_WORLD));
  //initializing NCCL, group API is required around ncclCommInitRank as it is
  //called across multiple GPUs in each thread/process
  NCCLCHECK(ncclGroupStart());
  for (int i=0; i<nDev; i++) {
     CUDACHECK(cudaSetDevice(localRank*nDev + i));
     NCCLCHECK(ncclCommInitRank(comms+i, nRanks*nDev, id, myRank*nDev + i));
  }
  NCCLCHECK(ncclGroupEnd());
  //calling NCCL communication API. Group API is required when using
  //multiple devices per thread/process
  NCCLCHECK(ncclGroupStart());
  for (int i=0; i<nDev; i++)
     NCCLCHECK(ncclAllReduce((const void*)sendbuff[i], (void*)recvbuff[i], size, ncclFloat, ncclSum,
           comms[i], s[i]));
  NCCLCHECK(ncclGroupEnd());
  //synchronizing on CUDA stream to complete NCCL communication
  for (int i=0; i<nDev; i++)
      CUDACHECK(cudaStreamSynchronize(s[i]));
  //freeing device memory
  for (int i=0; i<nDev; i++) {
     CUDACHECK(cudaFree(sendbuff[i]));
     CUDACHECK(cudaFree(recvbuff[i]));
  }
  //finalizing NCCL
  for (int i=0; i<nDev; i++) {
     ncclCommDestroy(comms[i]);
  }
  //finalizing MPI
  MPICHECK(MPI_Finalize());
  printf("[MPI Rank %d] Success \n", myRank);
  return 0;
}

这个例子里比较核心的部分在于：

1.建立communicator，作为一个通信任务
2.确定rank，也就是明确MPI的进程之间的拓扑关系
3.Rank 与 dveice的绑定，注意这里setdevice的操作完成了这一事业
4.以group 的方式提交collective communication的任务
5.进行异步的完成通知

这个例子还没有覆盖的地方在于对初始化过程没有涉及：

1.如何发现机器上的GPU和nvlink
2.如何initialize 机器上的GPU和nvlink、nvswitch
3.如何构建GPU之间的可通信的拓扑（明确可以通信的链路）
4.如何根据当前通信任务的通信数据量选择通信的nvlink数量
5.如何下发通信任务到硬件
6.计算任务和通信任务有关联吗，还是完全独立呢？

initilaize和异常的问题FM 一章会有说明，这一章着重回答通信任务相关的几个问题。

资源准备

建立拓扑

这一步getuniqueId 会调用。
第一步是探查当前机器上的所有可通信设备，比如Netwrok（NIC）、QPI（CPU）、PCIE（RC、PCIE switch）、nvlink（GPU，nvswitch），并根据各单元的active状态建立一张可通信图，此时通信图已经明确了两件事：有多少Node，有多少条path，及完成了确认可达性的任务。

拓扑计算

这一部分，在commRankInit 也会做。
第二步，为第一步发现的通信图，寻找一条各node 之间传输的最佳path，以带宽最大为基准。此处完成路径规划。
这里只是准备好可用的path，并不是真正通信就会占用这些数量的path/link，真正的通信是考虑利用率的，单任务占满带宽只会降低利用率，所以取得负载和所占path的平衡就是真正通信的关键。
CommRankInit
有了全局图，准备工作好了，怎么和具体任务挂钩呢？
Communicator 确定一个通信任务的覆盖范畴，也就是我这个通信任务包含哪些进程。
而我们之前提到，分布式通信topo是有结构的，或者是有序的，并不是无序而各点同质的，于是每个进程明确自己的rank，而对于使用GPU加速的进程，rank需要关联到具体GPU的ID。
所以当我们调用CommRankInit时，我们做了什么？将commID与rankId 配置给GPU。这样GPU用link发起通信时，也可以明白自己通信的范围，使用组播而不是广播方式增加不必要的带宽占用。

NCCL_API(ncclResult_t, ncclCommInitRank, ncclComm_t* newcomm, int nranks, ncclUniqueId commId, int myrank);
ncclResult_t ncclCommInitRank(ncclComm_t* newcomm, int nranks, ncclUniqueId commId, int myrank) {
  // Load the CUDA driver and dlsym hooks (can fail on old drivers)
  (void)ncclCudaLibraryInit();

  int cudaDev;
  ncclConfig_t config = NCCL_CONFIG_INITIALIZER;
  CUDACHECK(cudaGetDevice(&cudaDev));

  NvtxParamsCommInitRank payload{myrank, nranks, cudaDev};
  NVTX3_FUNC_WITH_PARAMS(CommInitRank, CommInitRankSchema, payload)

  NCCLCHECK(ncclCommInitRankDev(newcomm, nranks, commId, myrank, cudaDev, &config));
  return ncclSuccess;
}

最终这个CommInitRankDev也会作为一个async job 下给硬件（cudalaunch）的方式，完成配置。
问题来了，commid 和nranks是外部传入的参数，那么到底如何计算出来的呢？
id 的计算本身也是个问题，是在上层构建通信环的过程中所得。
CommInitRank 解决了逻辑上rank 和comm 对应的device，构建了comm的stream对象，所以nccl 每个communication有自己的stream。
有兴趣的同学可以自己追一下代码：

CommInitRankDev->CommInitRankFunc->CommAlloc->ncclStrongStreamConstruct
->cudaStreamCreateWithFlags,cudaEventCreateWithFlags

ComputeChannel

现在，我们已经获得了全局可通信的图，也知道了自己需要通信的范围，接下来，就可以根据具体的通信任务计算需要占用的link数量（静态带宽）了。
具体占用带宽的计算，我们略去不表，总体上看，通信任务本身最后也是cudaLaunchKernel，所以和一般cuda 任务一样，需要明确grid和block，channel 和 thread的选择最终就挂钩在grid和block上。
从目前代码的主体逻辑看，channel 计算有几个特点：

1.对于每个thread 可用传递的数据量有一个参考值，根据这个参考值明确当前任务需要多少thread
2.Channel 是 thread 上一层之概念，其实也就对应grid和block的关系
3.通信任务计算的过程，是寻找可以达到最优带宽的最少channel和thread 数量，先减少channel，当channel 为1，再减少threads，threads的计算单位是warp
4.除去负责主要通信任务的threads，还有用于管理和同步的warps（不同通信算法需要的管理warp数量不一，ring为1，tree为2）

最终计算出来的channel 数量，每一个channel 会有独立的launchkernel，一个channel 对应一个grid，channel 内的thread数量最后根据block的规格，计算block。每一个launch 对应一个plan，这个plan 里就是任务需要资源的相关信息，launch 本身就是提交plan。
另外，结合例子理解，所有的任务都可以被group 粒度去调度，也就是不是每一次 用户提交任务就立即下达至GPU，而是一个group 结束后异步下达至GPU的，group 调度粒度取决于配置和用户编程的调用。

吃果冻不吐果冻皮

专注于AI工程化（LLM、MLOps、LLMOps、RAG、Agent）落地。

148篇原创内容

公众号

enqueue调用流程

NCCLCommRankInit，完成资源准备后，就是下发通信指令了，下发通信API 如下：

NCCL_API(ncclResult_t, ncclReduce, const void* sendbuff, void* recvbuff, size_t count,
    ncclDataType_t datatype, ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream);
ncclResult_t ncclReduce(const void* sendbuff, void* recvbuff, size_t count,
    ncclDataType_t datatype, ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream) {
  struct NvtxParamsReduce {
    size_t bytes;
    int root;
    ncclRedOp_t op;
  };
  constexpr nvtxPayloadSchemaEntry_t ReduceSchema[] = {
    {0, NVTX_PAYLOAD_ENTRY_TYPE_SIZE, "Message size [bytes]"},
    {0, NVTX_PAYLOAD_ENTRY_TYPE_INT, "Root", nullptr, 0, offsetof(NvtxParamsReduce, root)},
    {0, NVTX_PAYLOAD_ENTRY_NCCL_REDOP, "Reduction operation", nullptr, 0,
      offsetof(NvtxParamsReduce, op)}
  };
  NvtxParamsReduce payload{count * ncclTypeSize(datatype), root, op};
  NVTX3_FUNC_WITH_PARAMS(Reduce, ReduceSchema, payload)

  struct ncclInfo info = { ncclFuncReduce, "Reduce",
    sendbuff, recvbuff, count, datatype, op, root, comm, stream, /* Args */
    REDUCE_CHUNKSTEPS, REDUCE_SLICESTEPS };
  return ncclEnqueueCheck(&info);
}

前面的代码主要是参数设置，核心在于ncclEnqueueCheck 。其基本逻辑如下：

host 这边主机异步提交job(包括通信命令），经过排队和group调度，通过cudalaunchkernel 发往GPU，GPU 内部MCU 跑device code。所以GPGPU的nvlink 实际收发逻辑在device 侧是毫无疑问的。

另外，不同的通信原语（reduce，gather，allreduce，scatter）有不同的通信接口，底下机制一致。

而不同类型的通信（peer2peer，collective）有不同的队列（queue），每个通信都会变成queue的task，每个task 都会被NCCL 编排成一个plan，plan 包括一次通信需要的资源（channel/grid数量，thread 数量/warps 对齐的）。

stream 机制

Stream 介绍过，作为host侧提交kernel 任务的组织单位。

Communication 自身的stream

  if (parent == NULL || !parent->config.splitShare) {
    struct ncclSharedResources* sharedRes = NULL;
    NCCLCHECK(ncclCalloc(&sharedRes, 1));
    /* most of attributes are assigned later in initTransportsRank(). */
    sharedRes->owner = comm;
    sharedRes->tpNRanks = comm->nRanks;
    NCCLCHECK(ncclCalloc(&sharedRes->tpRankToLocalRank, comm->nRanks));
    NCCLCHECK(ncclStrongStreamConstruct(&sharedRes->deviceStream));
    NCCLCHECK(ncclStrongStreamConstruct(&sharedRes->hostStream));
    comm->sharedRes = sharedRes;
    sharedRes->refCount = 1;
  } else {
    comm->sharedRes = parent->sharedRes;
    ncclAtomicRefCountIncrement(&parent->sharedRes->refCount);
  }

这里的hoststream与device stream 都是归属该commnucation 的stream，分别对应proxy 部分和nvlink 部分的kernel code。

任务的stream

我们看看ncclreduce的接口，看看info的传参吧。

NCCL_API(ncclResult_t, ncclReduce, const void* sendbuff, void* recvbuff, size_t count,
    ncclDataType_t datatype, ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream);

明明每一个communicator 本身就是拥有两个stream，为什么外部还需要传入一个stream呢？这个stream 是未知的，换句话说，极有可能是task所在的stream。
这里涉及到之前提到的cuda graph的能力。cuda graph 允许不同stream 之间协同工作。
Ncclreduce 一路调用，在taskappend，进行任务判断，会判断是否当前comm->task的stream和info->stream 一致，如果一致，说明本来就在一个stream里，如果不一致或者task为空，则comm->task 赋值为info stream。另一个值的注意的点是，假如一个comm 上执行了多次task，那么这些task 所在的stream 会维护一个list，目前看这个list是fifo的。

if (info->stream != tasks->streamRecent || tasks->streams == nullptr) {
    tasks->streamRecent = info->stream;
    struct ncclCudaStreamList* l = tasks->streams;
    while (true) {
      if (l == nullptr) { // Got to the end, this must be a new stream.
        struct ncclCudaGraph graph;
        NCCLCHECK(ncclCudaGetCapturingGraph(&graph, info->stream))
        if (tasks->streams != nullptr && !ncclCudaGraphSame(tasks->capturingGraph, graph)) {
          WARN("Streams given to a communicator within a NCCL group must either be all uncaptured or all captured by the same graph.");
          return ncclInvalidUsage;
        }
        tasks->capturingGraph = graph; // C++ struct assignment
        // Add stream to list
        l = ncclMemoryStackAlloc<struct ncclCudaStreamList>(&comm->memScoped);
        l->stream = info->stream;
        l->next = tasks->streams;
        tasks->streams = l;
        break;
      }
      if (l->stream == info->stream)
        break; // Already seen stream.
      l = l->next;
    }
  }

Launch 时的stream 关系

Info stream和comm的stream是什么关系呢？各stream 之间的关系在这里。
userstream【0】就是第一个用户task的stream，host和device的stream是comm的stream。

    // Semantically we want these dependencies for the kernels launched:
    //   1. Launch host task on hostStream.
    //   2. Launch kernel, depends on all of {deviceStream, hostStream, userStream[i]...}
    //   3. {deviceStream, userStream[i]...} depend on kernel.
    // We achieve this by:
    //   1. userStream[0] waits on deviceStream
    //   2. deviceStream waits on each of userStream[1...]
    //   3. host task launch on hostStream
    //   4. userStream[0] waits on hostStream
    //   5. kernel launch on userStream[0]
    //   6. deviceStream waits on userStream[0]
    //   7. userStream[1...] each waits on deviceStream
    // The two-level fan-in fan-out is because ncclStrongStreamWaitStream() requires
    // at least one of the two streams to be strong-stream.

在launch kernel 前，会尝试acquire comm的stream，host还是device 取决于task 类型。但是我们可以看到task stream和comm stream 存在互相依赖里了，并且说明了互相wait的关系。nvlink使用的都是devicestream ，rdma和tcp 使用的是hoststream。

计算通信之关系

细心的同学会在stream的逻辑中发现cudagraph的影子，我们知道cudagraph可以用来协同处理多个stream。
那么上述的graph究竟有没有计算通信协同呢？我们再过一遍这部分。
TaskAppend 时尝试获取stream 所在的graph，如果有graph，返回graph，否则返回空。

        NCCLCHECK(ncclCudaGetCapturingGraph(&graph, info->stream))

虽然该task 只是分布式通信任务，但是其stream所在的graph呢？不可知，毕竟是外部传入的，所以可能是计算任务的，或者更大的框架的。
而实际上ncclStrongStreamAcquire还是ncclStrongStreamWaitStream 都是允许graph 为0，所以即使task 的stream不在graph里，也不影响执行，逻辑只是削去了graph相关的部分，而只有event 的逻辑而已了。

ncclResult_t ncclStrongStreamAcquire(
    struct ncclCudaGraph graph, struct ncclStrongStream* ss
  ) {
  #if CUDART_VERSION >= 11030
    bool mixing = ncclParamGraphMixingSupport();
    if (graph.graph == nullptr) {
      if (mixing && ss->everCaptured) {
        CUDACHECK(cudaStreamWaitEvent(ss->cudaStream, ss->serialEvent, 0));
        ss->serialEventNeedsRecord = false;
      }
    } else {
      ss->everCaptured = true;
      // Find the current graph in our list of graphs if it exists.
      struct ncclStrongStreamGraph** pg = &ss->graphHead;
      struct ncclStrongStreamGraph* g;
      while (*pg != nullptr) {
        g = *pg;
        if (g->graphId == graph.graphId) {
          // Move to front of list so that operations after acquire don't have to search the list.
          *pg = g->next;
          g->next = ss->graphHead;
          ss->graphHead = g;
          return ncclSuccess;
        } else if (false == __atomic_load_n(&g->alive, __ATOMIC_ACQUIRE)) {
          // Unrelated graph that has been destroyed. Remove and delete.
          *pg = g->next;
          ncclStrongStreamGraphDelete(g);
        } else {
          pg = &g->next;
        }
      }

而如下的部分，实际上，如果task stream在图里，comm stream 也会被放在图里，一起执行；否则就是几条stream 互相event wait而已。device stream 等待第一个user stream（计算stream），其它user stream等待device stream上任务完成。

cudaStream_t launchStream = tasks->streams->stream;
NCCLCHECKGOTO(ncclStrongStreamAcquire(tasks->capturingGraph, &comm->sharedRes->deviceStream), result, failure);

// Create dependency for device stream on user streams. First from extra user
// streams to deviceStream. Then deviceStream to first user stream.
for (struct ncclCudaStreamList* l=tasks->streams->next; l != nullptr; l = l->next) {
    NCCLCHECKGOTO(ncclStrongStreamWaitStream(tasks->capturingGraph, &comm->sharedRes->deviceStream, l->stream), result, failure);
}

NCCLCHECKGOTO(ncclStrongStreamWaitStream(tasks->capturingGraph, launchStream, &comm->sharedRes->deviceStream), result, failure);

那么还有最后一个问题？nccl 最终会调用cudalaunch，但是graph 执行需要graph capture ending一起被graphlaunch，这里是怎么处理的呢？
如果存在图，当前task 会被认为是persistent 的task，会被拷贝至gpu上，在nccl clean group时也会绕过。后续流程就不得而知了。
从上述逻辑证明看，计算和通信是可以放在一个graph 内的，所以计算通信必然是可以协同的。这一部分可以参考系列三种cuda graph的部分说明。

Fabric manager 浅说

除去NCCL 这个众所周知的software stack，GPU 达到如此可用性，与Fabric manager 也是分不开的。

FM 是什么

FM 配置 NVSwitch memory fabric，以在所有参与的 GPU 之间形成一个内存结构，并监视支持该结构的 NVLink。在较高层面上，FM 具有以下职责：

1.配置NVSwitch 端口之间的路由。
2.与GPU驱动程序配合初始化GPU。
3.监控结构中的 NVLink 和 NVSwitch 错误。

在不支持基于自主链路初始化 (ALI) 的 NVLink 训练的系统（第一代和第二代基于 NVSwitch 的系统，即H100前）上，FM 还具有以下附加职责：

1.与 NVSwitch 驱动程序协调以训练 NVSwitch 到 NVSwitch NVLink 互连。
2.与 GPU 驱动程序协调，初始化并训练 NVSwitch 到 GPU NVLink 互连。

Nvlink 初始化

NVIDIA GPU 和 NVSwitch memory fabric是需要使用 NVIDIA 内核驱动程序的 PCIe 端点设备。在不支持 ALI 的 DGX-2、NVIDIA HGX-2、DGX A100 和 NVIDIA HGX A100 系统上，系统启动后，加载 NVIDIA 内核驱动程序后会启用 NVLink 连接，并且 FM 会配置这些连接。如果在 FM 完全初始化系统之前启动应用程序，或者 FM 无法初始化系统，则 CUDA 初始化将会失败，并出现 cudaErrorSystemNotReady 错误。在支持 ALI 的 DGX H100 和 NVIDIA HGX H100 系统上，NVLink 在 GPU 和 NVSwitch 硬件上进行训练并不需要FM。要启用 NVLink 对等支持，GPU 必须向 NVLink 结构注册。如果 GPU 无法注册到结构，它将失去其 NVLink 对等功能并可用于非对等用例。GPU 完成 NVLink 结构的注册过程后，CUDA 初始化过程将开始。

FM 模式

FM 主要配置四类模式：

1.初始化失败后的模式
2.Access link 失败的模式（access link指GPU和nvswitch间的link，nvlink）
3.Trunk link 失败的模式（trunk link 指switch 之间的link，osfp）
4.Switch 异常的模式（switch 本身失败）
5.FM 停止情况下job 的工作模式（GPU 上的任务失败）

具体失败类型请看下面的附录D
fabric-manager-user-guide.pdf

基本流程栈

这里看到从运维角度来说，完整的软件栈，包括NVML（nvidia-smi，monitor API）、GCGM（monitor backend-agent）、fabric manager service（backend service）、GPU & NVSWitch driver、BMC等等。

MIG 与 fabric manager

MIG 将 NVIDIA A100 或 H100 GPU 划分为许多独立的 GPU 实例。这些实例同时运行，每个实例都有自己的内存、缓存和流式多处理器。但是，当您启用 MIG 模式时，GPU NVLink 将被禁用，GPU 将失去其 NVLink 对等 (P2P) 功能。成功禁用 MIG 模式后，将再次启用 GPU NVLinks，恢复 GPU NVLink P2P 能力。在基于 NVSwitch 的 DGX 和 NVIDIA HGX 系统上，FM 服务可以与 GPU MIG 实例配合。此外，在这些系统上，要在禁用 MIG 模式后成功恢复 GPU NVLink 对等功能，FM 服务必须正在运行。在 DGX 上A100 和 NVIDIA HGX A100 系统，相应的 GPU NVLink 和 NVSwitch 侧 NVLink 在启用 MIG 模式时关闭，在禁用 MIG 模式时重新训练。但是，在 DGX H100 和 NVIDIA HGX H100 系统上，GPU NVLink 将在 MIG 模式期间保持active。

DCGM

NVIDIA 数据中心 GPU 管理器 (DCGM) 是一套用于在集群环境中管理和监控 NVIDIA 数据中心 GPU 的工具。它包括主动健康监控、全面诊断、系统警报和治理策略（包括电源和时钟管理），也包括link 状态、gpu 任务状态的监控。

https://https://docs.nvidia/datacenter/dcgm/1.6/pdf/dcgm-user-guide.pdf

NVML

管控工具，实际工具载体nvidia-smi。温度电压，板卡类型id等等，也包括GPU 利用率和活跃ce 数量。
https://developer.download.nvidia/assets/cuda/files/CUDADownloads/NVML/nvml.pdf
nvidia-smi topo -m 也可以用于获取当前机器的topo情况。

Host driver 的用处

目前看到的情况，Nvlink 和 nvswitch host driver 主要服务于fabric manager与nvml，与实际的数据面操作无关，以配置/获取管理信息为主，所以主体应用逻辑看fabric manager 这一部分即可。

Nvswitch

Nvswitch driver的功能相当淳朴，两部分：

1.nvswitch的初始化工作（probe）和去初始化
2.管理功能，主要包含以下几个类别：

• Nvswitch 信息获取，

  • 面向硬件，包括温度、bios、各种版本（driver、fw等）、内部延迟、硬件配置信息、计数器、状态、错误、i2c等等，

  • 面向访问，也包括ingress/egress request/response table，link table，link config等等

  • 面向流量，包数、带宽等等的统计信息

• 访问控制，比如修改routing 表，修改blacklist，直接读写寄存器

• 清理各种错误信息

nvlink

Nvlink 的功能相对复杂一些，不过总的而言也是几类：

1. GPU和nvlink 关系的绑定（add，remove）

2. nvlink的初始化和训练

3. Nvlink topo的发现，这一点主要是fabric manager和nvml 有获取topo和对端link信息的需求

4. Nvlink 连接情况

5. Nvlink 模式的变化，lower power state

主要是设备管理方面的能力，不过由于nvoc（一种nvidia内部语法糖，没看懂）的封装，这一部分看起来会比较模糊。

几种通信拓扑

Ring

如图所示，集群内的GPU 以ring的方式组织起来，假如有三张卡A，B，C，则通信过程为A->B,B->C,C->A。实际上ring 更多用于nvlink的情况。（考虑物理拓扑）

Tree

tree 与数据中心网络（fat tree）的物理拓扑比较接近，目前也主要用在数据中心网络中。

COLLNET-SHARP

sharp 是一种在网计算优化reduce 流程的架构，通过将reduce 计算offload到交换机上，降低了reduce 需要的通信负载和次数，有效带宽更高，latency更短。

历代网络架构

Cube-Mesh - 最简拓扑

link 均匀。

DGX-1 - 不均匀拓扑

注意V100 和 P100的差别，P100 cube上每条都是1 link，V100 cube上每个node 2条边2link，2条边1link。
在cube中的一个显著现象有的GPU访问需要2跳link，而有的只需要一跳，导致2跳的延迟高于1跳2倍。

这张图说明了为什么此类拓扑叫cube-mesh）。

DGX-2 - switch 出现

GTC2018发布的dgx-2，其加倍了v100的数量，最终高达16块v100。同时hbm2升级到32GB/块，一共高达512GB，cpu升级为双路2.7G 24核 xeon 8168.

dgx2拥有两块基板，这两块基板则是通过nvswitch剩余的另一侧接口完全互联在一起，这就构成了一个16路全连接的GPU构架。
两块基板之间的nvswitch之间都有八路link互联，16块GPU每块有6路nvlink的情况下，其总双路带宽达到2400GB/s。有趣的是，其实nvswitch有18路接口nvidia却只用到了其中16路。一种可能性是nv留下两路用于支持ibm的power9处理器（dgx1和2都是用的志强）。在这个复杂的结构中，power9处理器可能分别接在两块基板的nvsiwtch上，这样GPU也与Power9处于全连接状态。如果CPU直接与nvswitch相连，那么pcie就不再担任cpu与gpu相连的责任。目前nvidia还没有向其他厂商开放nvswitch，如果他们决定开放，将会产生一些新型态的，可能更加规模庞大的结算节点。

每个NVSwitch是18*18的crossbar交换机：

1. 8个port用于baseboard内的通信

2. 8个port用于和对端baseboard进行通信

3. 需要经过两跳交换机，但由于两边的交换机有一半是直连，所以可以当成一跳交换机

4. 另外两个port保留

DGX - A100

H100 superpod

Nv switch 收敛比 ingress：egress = 2:1。
一台机器有8张H100，4 个 nvswitch，每张H100 18link，以5、4、4、5 连接至4个 L1 nvswitch，每个nvswitch 单芯片，每芯片64link。
所以机器内的四个L1 nvswitch 承受连接数不同，每个link 50GB/s，所以四台L1 ingress分别是2 TB/s、1.6 TB/s、1.6 TB/s、2 TB/s，分别占用40、32、32、40 link。
L1 向上连接L2 switch，通过osfp，一个osfp 由4个link 组成，带宽为200Gbps。egress 方向2:1 收敛后分别为1TB/s，0.8TB/s，0.8TB/s，1TB/s，分别需要5、4、4、5条osfp，折合20、16、16、20 link。
L1 switch link 利用分别为（40 + 20）/64，（32+16）/64，（32+16）/64，（40+20）/64。
一个superpod 一共32台机器，32 * （1 + 0.8 + 0.8 + 1）TB/s = 115.2 TB/s = 32 * （5 + 4 + 4 + 5）osfp = 576 osfp，一台L2 16个osfp ，所以需要L2 switch 576/16 = 36台。

GH200 superpod

Grace hopper的拓扑示意图如下，依旧是两套网络，DPU 接quantum2 ，GPU 接 nvlink/nvswitch。

GH200 中 Nv switch 收敛比 ingress：egress = 1:1，相比H100 收敛比更高，意味总带宽可以提高。
一台机器中8张H100（带Grace CPU），每张H100 18link；3个switch，每个switch 2 芯片，每个芯片64link，每switch 合计128link。
每H100 以6、6、6 link连接至3台switch（每个switch chip 对应3link），比H100 均匀，每个chip 只有24link ingress，每个switch 48link ingress。
每个link 50GB/s，单switch chip 连接8*3=24link=1.2TB/s，单switch 2.4TB/s，1:1收敛比，所以egress 也是2.4TB/s，折合12 osfp，48link egress。
L1 的switch ingress + egress = 96link，占用率为96/128，32 port 空闲。
单台机器3个L1 switch，总egress 7.2TB/s = 36 osfp。
一个superPod 32 机器，L2 总ingress = 7.2TB/s * 32 = 36osfp * 32 = 1152 osfp。
L2 switch 规格与L1 一样，一台switch 128link = 32osfp，所以一共需要36台L2 switch。
虽然GH200 与 H100 在L2 switch 数量，机器数量以及GPU数量一致，但是由于收敛比设计不同，总带宽是不一样的，GH200 的一台switch 能力为H100 拓扑中的2倍。

除了NVlink 和 NVswitch 构成的high bandwidth networks, RDMA 构成了另一个网络，这整体构成railway架构。我们之前讲述的是nvlink 网络的架构，侧重pod内互联，pod内GPU通信通过nvlink，而pod间通过RDMA。RDMA 网络的结构如下图，假设我们有M个pod，每个pod内有K张GPU，pod内K个GPU通过nvlink和nvswitch互联，每个pod内的第i个GPU通过RDMA连接到同一个switch上，所以一共有K个rail switch，因此不同pod之间可以通过 rail switch互联，而rail switch都连接到spine switch，所以整个网络可以互联。spine switch 可以看作是对外的接口，以及rail switch 不能满足需求时的保障。然而实际上，大模型训练过程中nvlink 占通信70%以上，其余主要集中在rail switch，spine switch上很少流量。

附录1- NVprof / profling 工具

Nvidia profling 工具，包括GPU、link、CPU 上的性能数据抓取，对于nvlink来说主要是各种带宽数据和topo。以下仅举例：

Visual Profiler collects NVLink topology and NVLink transmit/receive throughput metrics and maps the metrics on to the topology. The topology is collected by default along with the timeline. Throughput/ utilization metrics are generated only when NVLink option is chosen. NVLink information is presented in the Results section of Examine GPU Usage in CUDA Application Analysis in Guided Analysis. NVLink Analysis shows topology that shows the logical NVLink connections between different devices. A logical link comprises of 1 to 4 physical NVLinks of same properties connected between two devices. Visual profiler lists the properties and achieved utilization for logical NVLinks in ‘Logical NVLink Properties’ table. It also lists the transmit and receive throughputs for logical NVLink in ‘Logical NVLink Throughput’ table.

官方文档在此：https://docs.nvidia/cuda/pdf/CUDA_Profiler_Users_Guide.pdf

附录2-PXN - PCI X NVLINK

The new feature introduced in NCCL 2.12 is called PXN, as PCI × NVLink, as it enables a GPU to communicate with a NIC on the node through NVLink and then PCI. This is instead of going through the CPU using QPI or other inter-CPU protocols, which would not be able to deliver full bandwidth. That way, even though each GPU still tries to use its local NIC as much as possible, it can reach other NICs if required. 就是机外nic，机内nvlink，虽然是个新feature，但用的老拓扑。不过机外网络训练使用的doubling的实现。

本文标签： 架构通信NVIDIAGPGPU