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

摆烂了整整一年。
工作的事情,真的影响心情。

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

NCCL
集合通讯
P2P
CC
Broadcast\Reduce\Gather...
算法和网络拓扑结构
协议
RING
TREE
COLLNET
NCCL_PROTO_SIMPLE
NCCL_PROTO_LL
NCCL_PROTO_LL128

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

背景

是什么限制了并行计算应用?
1 并行计算任务的效率

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

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

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

集合通讯

任务之间有哪些通讯模式
Point-to-point communication (P2P, 也就是点对点)

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

Collective communication(CC, 集合通信)

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

P2P,点对点通信

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

CC, 集合通信

cc

Broadcast, 广播

BROADCAST

Scatter,单发多收

scatter

Gather,多发单收

Gather

All Gather

All Gather

Reduce

Reduce

All Reduce

All Reduce

Reduce-Scatter

Reduce-Scatter

All to All

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=<path to cuda install>

直接编译:

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

  • 7
    点赞
  • 23
    收藏
    觉得还不错? 一键收藏
  • 打赏
    打赏
  • 1
    评论

“相关推荐”对你有帮助么?

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论 1
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包

打赏作者

canmoumou

你的鼓励将是我创作的最大动力

¥1 ¥2 ¥4 ¥6 ¥10 ¥20
扫码支付:¥1
获取中
扫码支付

您的余额不足,请更换扫码支付或充值

打赏作者

实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

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

余额充值