【分布式】NCCL部署与测试 - 01
作者:mmseoamin日期:2023-12-18

摆烂了整整一年。

工作的事情,真的影响心情。

目录

  • NCCL
    • 相关系列
    • 简述
    • 背景
    • 集合通讯
      • P2P,点对点通信
      • CC, 集合通信
        • Broadcast, 广播
        • Scatter,单发多收
        • Gather,多发单收
        • All Gather
        • Reduce
        • All Reduce
        • Reduce-Scatter
        • All to All
        • 可能存在的问题
        • 代码结构
          • 编译
          • 测试
          • 其他
            • 1、Group
            • 2、Sendrecv
            • 相关系列

              NCCL

              相关系列

              分布式】NCCL部署与测试 - 01

              【分布式】入门级NCCL多机并行实践 - 02

              【分布式】小白看Ring算法 - 03

              【分布式】大模型分布式训练入门与实践 - 04

              简述

              NCCL(NVIDIA Collective Communications Library)是一种由NVIDIA开发的高性能通信库,专门用于加速多GPU系统中的并行计算工作负载。NCCL旨在优化多GPU之间的数据传输和通信,以实现更快的并行计算速度。

              也就是说NCCL可以加速GPU通信,降低通信开销

              NCCL具备以下通信模式(后文有详细介绍):

              点对点通信:允许两个特定的GPU之间直接交换数据。

              群集通信:NCCL还支持集体通信,这些操作涉及多个GPU之间的数据交换

              NCCL还具有的其他功能

              拓扑感知通信:识别多GPU系统的拓扑结构,并优化通信以最大程度地减少通信延迟和带宽消耗。

              异步通信:允许计算操作与通信操作并行执行,从而提高了系统的效率。

              那么也就是说我们可以通过一下几个方向来了解什么是NCCL

              本章节只放了集合通讯相关和一些零碎的内容,后续再讲其他的。

              背景

              是什么限制了并行计算应用?

              1 并行计算任务的效率

              • 可供使用的并行数量
              • 分配给每个处理器的任务

                2 各个任务之间的通信开销

                • 通信量
                • 计算与通信的重叠程度

                  集合通讯

                  任务之间有哪些通讯模式

                  Point-to-point communication (P2P, 也就是点对点)

                  • 单一发送方,单一接收方
                  • 高效通信相对好实现

                    Collective communication(CC, 集合通信)

                    • 多个发送方和/或接收方
                    • 模式有broadcast, scatter, gather, reduce, all-to-all, …
                    • 很难做到高效

                      P2P,点对点通信

                      P2P是HPC 中最常见的模式,其中通信通常是与最近的邻居

                      【分布式】NCCL部署与测试 - 01,P2P,第1张

                      CC, 集合通信

                      【分布式】NCCL部署与测试 - 01,cc,第2张

                      Broadcast, 广播

                      【分布式】NCCL部署与测试 - 01,BROADCAST,第3张

                      Scatter,单发多收

                      【分布式】NCCL部署与测试 - 01,scatter,第4张

                      Gather,多发单收

                      【分布式】NCCL部署与测试 - 01,Gather,第5张

                      All Gather

                      【分布式】NCCL部署与测试 - 01,All Gather,第6张

                      Reduce

                      【分布式】NCCL部署与测试 - 01,Reduce,第7张

                      All Reduce

                      【分布式】NCCL部署与测试 - 01,All Reduce,第8张

                      Reduce-Scatter

                      【分布式】NCCL部署与测试 - 01,Reduce-Scatter,第9张

                      All to All

                      【分布式】NCCL部署与测试 - 01,All to All,第10张

                      可能存在的问题

                      1. 有多个发送方和/或接收方会加剧通信效率低下

                        • 对于小额传输,延迟占主导地位,更多参与者会增加延迟

                        • 对于大型传输,带宽是关键,瓶颈很容易暴露出来

                        • 可能需要拓扑感知实现以实现高性能

                        • CC通常是阻塞/不重叠的

                      2. 使用场景

                        • 深度学习 (All-reduce, broadcast, gather)

                        • 并行FFT 傅里叶变换(Transposition is all-to-all)

                        • 分子动力学 (All-reduce)

                        • 图形分析(All-to-all)

                      3. 如何扩展?需要高效的通信算法和谨慎的实现

                        • 通信算法依赖于拓扑

                        • 拓扑可能很复杂 - 并非每个系统都是一棵肥树

                        • 大多数集合体都适合在环上实现带宽优化,并且许多拓扑可以解释为一个或多个环[P. Patarasuk and X. Yuan]

                      代码结构

                      仓库:NCCL Github

                      测试仓库:Nccl-test Github

                      编译

                      可直接参考Readme。

                      默认有cuda环境不用设置,如果有多个cuda环境,需要在make的时候指定CUDA_HOME。

                      git clone https://github.com/NVIDIA/nccl.git
                      cd nccl
                      make src.build CUDA_HOME=
                      

                      直接编译:

                      make -j src.build
                      

                      测试

                      运行时可以从nccl-test入手。同样参考Github的README即可。

                      NCCL和nccl-test都是使用makefile编写的链接。

                      git https://github.com/NVIDIA/nccl-tests.git
                      cd nccl-tests
                      make CUDA_HOME=/path/to/cuda NCCL_HOME=/path/to/nccl
                      

                      测试时:

                      Run on 8 GPUs (-g 8), scanning from 8 Bytes to 128MBytes :

                      $ ./build/all_reduce_perf -b 8 -e 128M -f 2 -g 8
                      

                      Run with MPI on 10 processes (potentially on multiple nodes) with 4 GPUs each, for a total of 40 GPUs:

                      $ mpirun -np 10 ./build/all_reduce_perf -b 8 -e 128M -f 2 -g 4
                      

                      其他

                      1、Group

                      通过任意一个nccl-test,我们可以看到大部分原语操作都要在前后包裹ncclGroupStart()和ncclGroupEnd()。这个搭配是成对的,并且可以嵌套完成。Start开始以后,会根据你调用接口的同步异步模式进行差异处理,最后调用End才是真正执行。

                      如果是同步模式,则Start以后会立即调用,如果是异步模式,则会直接把这一组操作加入任务队列。

                      NCCL_API(ncclResult_t, ncclGroupStart);
                      ncclResult_t ncclGroupStart() {
                        ncclResult_t ret = ncclSuccess;
                        NVTX3_FUNC_RANGE_IN(nccl_domain);
                        NCCLCHECK(ncclGroupStartInternal());
                        TRACE_CALL("ncclGroupStart()");
                        return ret;
                      }
                      NCCL_API(ncclResult_t, ncclGroupEnd);
                      ncclResult_t ncclGroupEnd() {
                        ncclResult_t ret = ncclSuccess;
                        NVTX3_FUNC_RANGE_IN(nccl_domain);
                        NCCLCHECKGOTO(ncclGroupEndInternal(), ret, exit);
                        TRACE_CALL("ncclGroupEnd()");
                      exit:
                        return ret;
                        }
                      

                      2、Sendrecv

                      另外,代码中实际存在的原语只有:

                      (在src/collectives/目录下)