【分布式】NCCL部署与测试 – 01

摆烂了整整一年。

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

目录

• 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 中最常见的模式，其中通信通常是与最近的邻居

CC, 集合通信

Broadcast, 广播

Scatter，单发多收

Gather，多发单收

All Gather

Reduce

All Reduce

Reduce-Scatter

All to All

可能存在的问题

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/目录下）

• all_gather.cc
• all_reduce.cc
• broadcast.cc
• reduce.cc
• reduce_scatter.cc
• sendrecv.cc

也就是说，实际应用场景下，其他的模式全都是用sendrecv.cc实现的，包括all-to-all等。

相关系列

【分布式】NCCL部署与测试 – 01

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

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

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