摆烂了整整一年。
工作的事情,真的影响心情。
目录
- 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
#mermaid-svg-2e6JoUvFSvCGo1px {font-family:”trebuchet ms”,verdana,arial,sans-serif;font-size:16px;fill:#333;}#mermaid-svg-2e6JoUvFSvCGo1px .error-icon{fill:#552222;}#mermaid-svg-2e6JoUvFSvCGo1px .error-text{fill:#552222;stroke:#552222;}#mermaid-svg-2e6JoUvFSvCGo1px .edge-thickness-normal{stroke-width:2px;}#mermaid-svg-2e6JoUvFSvCGo1px .edge-thickness-thick{stroke-width:3.5px;}#mermaid-svg-2e6JoUvFSvCGo1px .edge-pattern-solid{stroke-dasharray:0;}#mermaid-svg-2e6JoUvFSvCGo1px .edge-pattern-dashed{stroke-dasharray:3;}#mermaid-svg-2e6JoUvFSvCGo1px .edge-pattern-dotted{stroke-dasharray:2;}#mermaid-svg-2e6JoUvFSvCGo1px .marker{fill:#333333;stroke:#333333;}#mermaid-svg-2e6JoUvFSvCGo1px .marker.cross{stroke:#333333;}#mermaid-svg-2e6JoUvFSvCGo1px svg{font-family:”trebuchet ms”,verdana,arial,sans-serif;font-size:16px;}#mermaid-svg-2e6JoUvFSvCGo1px .label{font-family:”trebuchet ms”,verdana,arial,sans-serif;color:#333;}#mermaid-svg-2e6JoUvFSvCGo1px .cluster-label text{fill:#333;}#mermaid-svg-2e6JoUvFSvCGo1px .cluster-label span{color:#333;}#mermaid-svg-2e6JoUvFSvCGo1px .label text,#mermaid-svg-2e6JoUvFSvCGo1px span{fill:#333;color:#333;}#mermaid-svg-2e6JoUvFSvCGo1px .node rect,#mermaid-svg-2e6JoUvFSvCGo1px .node circle,#mermaid-svg-2e6JoUvFSvCGo1px .node ellipse,#mermaid-svg-2e6JoUvFSvCGo1px .node polygon,#mermaid-svg-2e6JoUvFSvCGo1px .node path{fill:#ECECFF;stroke:#9370DB;stroke-width:1px;}#mermaid-svg-2e6JoUvFSvCGo1px .node .label{text-align:center;}#mermaid-svg-2e6JoUvFSvCGo1px .node.clickable{cursor:pointer;}#mermaid-svg-2e6JoUvFSvCGo1px .arrowheadPath{fill:#333333;}#mermaid-svg-2e6JoUvFSvCGo1px .edgePath .path{stroke:#333333;stroke-width:2.0px;}#mermaid-svg-2e6JoUvFSvCGo1px .flowchart-link{stroke:#333333;fill:none;}#mermaid-svg-2e6JoUvFSvCGo1px .edgeLabel{background-color:#e8e8e8;text-align:center;}#mermaid-svg-2e6JoUvFSvCGo1px .edgeLabel rect{opacity:0.5;background-color:#e8e8e8;fill:#e8e8e8;}#mermaid-svg-2e6JoUvFSvCGo1px .cluster rect{fill:#ffffde;stroke:#aaaa33;stroke-width:1px;}#mermaid-svg-2e6JoUvFSvCGo1px .cluster text{fill:#333;}#mermaid-svg-2e6JoUvFSvCGo1px .cluster span{color:#333;}#mermaid-svg-2e6JoUvFSvCGo1px div.mermaidTooltip{position:absolute;text-align:center;max-width:200px;padding:2px;font-family:”trebuchet ms”,verdana,arial,sans-serif;font-size:12px;background:hsl(80, 100%, 96.2745098039%);border:1px solid #aaaa33;border-radius:2px;pointer-events:none;z-index:100;}#mermaid-svg-2e6JoUvFSvCGo1px :root{–mermaid-font-family:”trebuchet ms”,verdana,arial,sans-serif;}
本章节只放了集合通讯相关和一些零碎的内容,后续再讲其他的。
背景
是什么限制了并行计算应用?
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
可能存在的问题
- 有多个发送方和/或接收方会加剧通信效率低下
• 对于小额传输,延迟占主导地位,更多参与者会增加延迟
• 对于大型传输,带宽是关键,瓶颈很容易暴露出来
• 可能需要拓扑感知实现以实现高性能
• CC通常是阻塞/不重叠的 - 使用场景
• 深度学习 (All-reduce, broadcast, gather)
• 并行FFT 傅里叶变换(Transposition is all-to-all)
• 分子动力学 (All-reduce)
• 图形分析(All-to-all) - 如何扩展?需要高效的通信算法和谨慎的实现
• 通信算法依赖于拓扑
• 拓扑可能很复杂 – 并非每个系统都是一棵肥树
• 大多数集合体都适合在环上实现带宽优化,并且许多拓扑可以解释为一个或多个环[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