ToB企服应用市场:ToB评测及商务社交产业平台

标题: NVIDIA GPGPU通讯架构 [打印本页]

---

作者: 怀念夏天    时间: 2024-12-14 00:53
标题: NVIDIA GPGPU通讯架构
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的新特性包括：

• 更高的带宽、信号线数减半、单GPU 12个nvlink
  
• 升级error detection和recovery。
  
• Write利用变成non-posted，使得请求侧可以举行同步，错误处理也有改进
  
• 优化了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.com/mp/appmsgalbum
oneflow 的剖析写得十分具体（抱拳），我这里仅做一点本身读源码的总结。
我们这里不讲一般NIC network相关的软件栈。主要会合于GPU和nvlink的链路。
基本例子

简单的聚集通讯例子

1. int main(int argc, char* argv[])
2. {
3.   int size = 32*1024*1024;
4.   int myRank, nRanks, localRank = 0;
5.   //initializing MPI
6.   MPICHECK(MPI_Init(&argc, &argv));
7.   MPICHECK(MPI_Comm_rank(MPI_COMM_WORLD, &myRank));
8.   MPICHECK(MPI_Comm_size(MPI_COMM_WORLD, &nRanks));
9.   //calculating localRank which is used in selecting a GPU
10.   uint64_t hostHashs[nRanks];
11.   char hostname[1024];
12.   getHostName(hostname, 1024);
13.   hostHashs[myRank] = getHostHash(hostname);
14.   MPICHECK(MPI_Allgather(MPI_IN_PLACE, 0, MPI_DATATYPE_NULL, hostHashs, sizeof(uint64_t), MPI_BYTE, MPI_COMM_WORLD));
15.   for (int p=0; p<nRanks; p++) {
16.      if (p == myRank) break;
17.      if (hostHashs[p] == hostHashs[myRank]) localRank++;
18.   }
19.   //each process is using two GPUs
20.   int nDev = 2;
21.   float** sendbuff = (float**)malloc(nDev * sizeof(float*));
22.   float** recvbuff = (float**)malloc(nDev * sizeof(float*));
23.   cudaStream_t* s = (cudaStream_t*)malloc(sizeof(cudaStream_t)*nDev);
24.   //picking GPUs based on localRank
25.   for (int i = 0; i < nDev; ++i) {
26.     CUDACHECK(cudaSetDevice(localRank*nDev + i));
27.     CUDACHECK(cudaMalloc(sendbuff + i, size * sizeof(float)));
28.     CUDACHECK(cudaMalloc(recvbuff + i, size * sizeof(float)));
29.     CUDACHECK(cudaMemset(sendbuff[i], 1, size * sizeof(float)));
30.     CUDACHECK(cudaMemset(recvbuff[i], 0, size * sizeof(float)));
31.     CUDACHECK(cudaStreamCreate(s+i));
32.   }
33.   ncclUniqueId id;
34.   ncclComm_t comms[nDev];
35.   //generating NCCL unique ID at one process and broadcasting it to all
36.   if (myRank == 0) ncclGetUniqueId(&id);
37.   MPICHECK(MPI_Bcast((void *)&id, sizeof(id), MPI_BYTE, 0, MPI_COMM_WORLD));
38.   //initializing NCCL, group API is required around ncclCommInitRank as it is
39.   //called across multiple GPUs in each thread/process
40.   NCCLCHECK(ncclGroupStart());
41.   for (int i=0; i<nDev; i++) {
42.      CUDACHECK(cudaSetDevice(localRank*nDev + i));
43.      NCCLCHECK(ncclCommInitRank(comms+i, nRanks*nDev, id, myRank*nDev + i));
44.   }
45.   NCCLCHECK(ncclGroupEnd());
46.   //calling NCCL communication API. Group API is required when using
47.   //multiple devices per thread/process
48.   NCCLCHECK(ncclGroupStart());
49.   for (int i=0; i<nDev; i++)
50.      NCCLCHECK(ncclAllReduce((const void*)sendbuff[i], (void*)recvbuff[i], size, ncclFloat, ncclSum,
51.            comms[i], s[i]));
52.   NCCLCHECK(ncclGroupEnd());
53.   //synchronizing on CUDA stream to complete NCCL communication
54.   for (int i=0; i<nDev; i++)
55.       CUDACHECK(cudaStreamSynchronize(s[i]));
56.   //freeing device memory
57.   for (int i=0; i<nDev; i++) {
58.      CUDACHECK(cudaFree(sendbuff[i]));
59.      CUDACHECK(cudaFree(recvbuff[i]));
60.   }
61.   //finalizing NCCL
62.   for (int i=0; i<nDev; i++) {
63.      ncclCommDestroy(comms[i]);
64.   }
65.   //finalizing MPI
66.   MPICHECK(MPI_Finalize());
67.   printf("[MPI Rank %d] Success \n", myRank);
68.   return 0;
69. }

复制代码

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

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发起通讯时，也可以明白本身通讯的范围，使用组播而不是广播方式增长不须要的带宽占用。

1. NCCL_API(ncclResult_t, ncclCommInitRank, ncclComm_t* newcomm, int nranks, ncclUniqueId commId, int myrank);
2. ncclResult_t ncclCommInitRank(ncclComm_t* newcomm, int nranks, ncclUniqueId commId, int myrank) {
3.   // Load the CUDA driver and dlsym hooks (can fail on old drivers)
4.   (void)ncclCudaLibraryInit();
6.   int cudaDev;
7.   ncclConfig_t config = NCCL_CONFIG_INITIALIZER;
8.   CUDACHECK(cudaGetDevice(&cudaDev));
10.   NvtxParamsCommInitRank payload{myrank, nranks, cudaDev};
11.   NVTX3_FUNC_WITH_PARAMS(CommInitRank, CommInitRankSchema, payload)
13.   NCCLCHECK(ncclCommInitRankDev(newcomm, nranks, commId, myrank, cudaDev, &config));
14.   return ncclSuccess;
15. }

复制代码

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

1. CommInitRankDev->CommInitRankFunc->CommAlloc->ncclStrongStreamConstruct
2. ->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 如下：

1. NCCL_API(ncclResult_t, ncclReduce, const void* sendbuff, void* recvbuff, size_t count,
2.     ncclDataType_t datatype, ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream);
3. ncclResult_t ncclReduce(const void* sendbuff, void* recvbuff, size_t count,
4.     ncclDataType_t datatype, ncclRedOp_t op, int root, ncclComm_t comm, cudaStream_t stream) {
5.   struct NvtxParamsReduce {
6.     size_t bytes;
7.     int root;
8.     ncclRedOp_t op;
9.   };
10.   constexpr nvtxPayloadSchemaEntry_t ReduceSchema[] = {
11.     {0, NVTX_PAYLOAD_ENTRY_TYPE_SIZE, "Message size [bytes]"},
12.     {0, NVTX_PAYLOAD_ENTRY_TYPE_INT, "Root", nullptr, 0, offsetof(NvtxParamsReduce, root)},
13.     {0, NVTX_PAYLOAD_ENTRY_NCCL_REDOP, "Reduction operation", nullptr, 0,
14.       offsetof(NvtxParamsReduce, op)}
15.   };
16.   NvtxParamsReduce payload{count * ncclTypeSize(datatype), root, op};
17.   NVTX3_FUNC_WITH_PARAMS(Reduce, ReduceSchema, payload)
19.   struct ncclInfo info = { ncclFuncReduce, "Reduce",
20.     sendbuff, recvbuff, count, datatype, op, root, comm, stream, /* Args */
21.     REDUCE_CHUNKSTEPS, REDUCE_SLICESTEPS };
22.   return ncclEnqueueCheck(&info);
23. }

复制代码

前面的代码主要是参数设置，核心在于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

1.   if (parent == NULL || !parent->config.splitShare) {
2.     struct ncclSharedResources* sharedRes = NULL;
3.     NCCLCHECK(ncclCalloc(&sharedRes, 1));
4.     /* most of attributes are assigned later in initTransportsRank(). */
5.     sharedRes->owner = comm;
6.     sharedRes->tpNRanks = comm->nRanks;
7.     NCCLCHECK(ncclCalloc(&sharedRes->tpRankToLocalRank, comm->nRanks));
8.     NCCLCHECK(ncclStrongStreamConstruct(&sharedRes->deviceStream));
9.     NCCLCHECK(ncclStrongStreamConstruct(&sharedRes->hostStream));
10.     comm->sharedRes = sharedRes;
11.     sharedRes->refCount = 1;
12.   } else {
13.     comm->sharedRes = parent->sharedRes;
14.     ncclAtomicRefCountIncrement(&parent->sharedRes->refCount);
15.   }

复制代码

这里的hoststream与device stream 都是归属该commnucation 的stream，分别对应proxy 部分和nvlink 部分的kernel code。
任务的stream
我们看看ncclreduce的接口，看看info的传参吧。

1. NCCL_API(ncclResult_t, ncclReduce, const void* sendbuff, void* recvbuff, size_t count,
2.     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的。

1. if (info->stream != tasks->streamRecent || tasks->streams == nullptr) {
2.     tasks->streamRecent = info->stream;
3.     struct ncclCudaStreamList* l = tasks->streams;
4.     while (true) {
5.       if (l == nullptr) { // Got to the end, this must be a new stream.
6.         struct ncclCudaGraph graph;
7.         NCCLCHECK(ncclCudaGetCapturingGraph(&graph, info->stream))
8.         if (tasks->streams != nullptr && !ncclCudaGraphSame(tasks->capturingGraph, graph)) {
9.           WARN("Streams given to a communicator within a NCCL group must either be all uncaptured or all captured by the same graph.");
10.           return ncclInvalidUsage;
11.         }
12.         tasks->capturingGraph = graph; // C++ struct assignment
13.         // Add stream to list
14.         l = ncclMemoryStackAlloc<struct ncclCudaStreamList>(&comm->memScoped);
15.         l->stream = info->stream;
16.         l->next = tasks->streams;
17.         tasks->streams = l;
18.         break;
19.       }
20.       if (l->stream == info->stream)
21.         break; // Already seen stream.
22.       l = l->next;
23.     }
24.   }

复制代码

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

1.     // Semantically we want these dependencies for the kernels launched:
2.     //   1. Launch host task on hostStream.
3.     //   2. Launch kernel, depends on all of {deviceStream, hostStream, userStream[i]...}
4.     //   3. {deviceStream, userStream[i]...} depend on kernel.
5.     // We achieve this by:
6.     //   1. userStream[0] waits on deviceStream
7.     //   2. deviceStream waits on each of userStream[1...]
8.     //   3. host task launch on hostStream
9.     //   4. userStream[0] waits on hostStream
10.     //   5. kernel launch on userStream[0]
11.     //   6. deviceStream waits on userStream[0]
12.     //   7. userStream[1...] each waits on deviceStream
13.     // The two-level fan-in fan-out is because ncclStrongStreamWaitStream() requires
14.     // 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，否则返回空。

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

复制代码

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

1. ncclResult_t ncclStrongStreamAcquire(
2.     struct ncclCudaGraph graph, struct ncclStrongStream* ss
3.   ) {
4.   #if CUDART_VERSION >= 11030
5.     bool mixing = ncclParamGraphMixingSupport();
6.     if (graph.graph == nullptr) {
7.       if (mixing && ss->everCaptured) {
8.         CUDACHECK(cudaStreamWaitEvent(ss->cudaStream, ss->serialEvent, 0));
9.         ss->serialEventNeedsRecord = false;
10.       }
11.     } else {
12.       ss->everCaptured = true;
13.       // Find the current graph in our list of graphs if it exists.
14.       struct ncclStrongStreamGraph** pg = &ss->graphHead;
15.       struct ncclStrongStreamGraph* g;
16.       while (*pg != nullptr) {
17.         g = *pg;
18.         if (g->graphId == graph.graphId) {
19.           // Move to front of list so that operations after acquire don't have to search the list.
20.           *pg = g->next;
21.           g->next = ss->graphHead;
22.           ss->graphHead = g;
23.           return ncclSuccess;
24.         } else if (false == __atomic_load_n(&g->alive, __ATOMIC_ACQUIRE)) {
25.           // Unrelated graph that has been destroyed. Remove and delete.
26.           *pg = g->next;
27.           ncclStrongStreamGraphDelete(g);
28.         } else {
29.           pg = &g->next;
30.         }
31.       }

复制代码

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

1. cudaStream_t launchStream = tasks->streams->stream;
2. NCCLCHECKGOTO(ncclStrongStreamAcquire(tasks->capturingGraph, &comm->sharedRes->deviceStream), result, failure);
4. // Create dependency for device stream on user streams. First from extra user
5. // streams to deviceStream. Then deviceStream to first user stream.
6. for (struct ncclCudaStreamList* l=tasks->streams->next; l != nullptr; l = l->next) {
7.     NCCLCHECKGOTO(ncclStrongStreamWaitStream(tasks->capturingGraph, &comm->sharedRes->deviceStream, l->stream), result, failure);
8. }
10. 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.com/datacenter/dcgm/1.6/pdf/dcgm-user-guide.pdf
NVML

管控工具，现实工具载体nvidia-smi。温度电压，板卡类型id等等，也包括GPU 使用率和活跃ce 数量。
https://developer.download.nvidia.cn/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 的功能相对复杂一些，不过总的而言也是几类：

• GPU和nvlink 关系的绑定（add，remove）
  
• nvlink的初始化和训练
  
• Nvlink topo的发现，这一点主要是fabric manager和nvml 有获取topo和对端link信息的需求
  
• Nvlink 连接情况
  
• 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交换机：

• 8个port用于baseboard内的通讯
  
• 8个port用于和对端baseboard举行通讯
  
• 必要经过两跳交换机，但由于两边的交换机有一半是直连，以是可以当成一跳交换机
  
• 别的两个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.com/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的实现。

免责声明：如果侵犯了您的权益，请联系站长，我们会及时删除侵权内容，谢谢合作！更多信息从访问主页：qidao123.com:ToB企服之家，中国第一个企服评测及商务社交产业平台。

---

欢迎光临 ToB企服应用市场:ToB评测及商务社交产业平台 (https://dis.qidao123.com/)

Powered by Discuz! X3.4