H100 · 多 GPU 编程（Multi GPU Part 2）

目录 · ← l11

H100 · 多 GPU 编程（Multi GPU Part 2）

本文档基于课程讲义《10. Multi GPU Part 2.pdf》（Lesson 10，60 页）整理， 系统介绍多 GPU 分布式训练的作业调度（Slurm/PMIx）、集合通信（NCCL）、四种并行策略。 前置阅读：H100-Multi-GPU.md（硬件互联体系）。

目录

1. Slurm 与 PMIx（作业调度与进程发现）
2. 设备隔离与 GPU 绑定
3. NCCL 概述与初始化
4. NCCL 六大集合通信原语
5. AI 的四种并行策略
6. NCCL 各操作详解
7. 点对点与分组：ncclSend/Recv 与 ncclGroup
8. 总结与学习衔接

1. Slurm 与 PMIx（作业调度与进程发现）

1.1 Slurm

• Slurm 是开源的作业管理器/调度器，面向 Linux 高性能计算集群。
• 它分配 GPU、CPU、内存；知道你的作业在哪跑，但不一定知道你的应用如何跨节点自通信。
• 过去 mpirun 用 SSH 连每台节点、核对主机名、交换密钥；大集群（64+ 节点）这个”握手”可能要好几分钟。
• 现在用 srun --mpi=pmix：Slurm 同时在所有节点上启动进程。

1.2 PMIx（Exascale 进程管理接口）

• Slurm 启动 1000 个进程后，rank 5 知道自己活着，但不知道 rank 0 的 IP、也不知道怎么联系 rank 999——它们是孤岛。
• 训练开始前，进程需要交换技术细节（IP、GPU 句柄）。PMIx 提供一个临时数据库：
  • PMIx_Put：进程把元数据（主机 IP、端口、CUDA IPC 句柄）推到本地 PMIx server。
  • PMIx_Commit：把本地数据推到全局命名空间。
  • PMIx_Get：其它进程查询这些数据以发现对等方。

1.3 为什么 Slurm + PMIx 一起重要

• Slurm 单独负责资源分配（节点/GPU/CPU）与任务放置，但历史上依赖旧 PMI（扩展性难超 ~1 万 GPU）。
• PMIx 是现代、面向 exascale 的进程管理接口（取代 PMI-1/PMI-2）。
• 二者集成带来：许多情况下免 mpirun 直接启动、大规模高效交换作业信息（rank/节点列表/端点）、 通过 NCCL-PMIx 插件与 NCCL 紧密集成。

1.4 Slurm 脚本与术语

• 脚本（.sh/.sbatch）告诉 Slurm 两件事：要什么资源（时间/GPU/CPU/内存）、拿到后做什么（跑 Python、编译等）。
• 用 sbatch script.sh 提交；Slurm 登录计算节点、设置环境、运行脚本命令。

2. 设备隔离与 GPU 绑定

2.1 设备隔离（两个机制）

若让 Slurm 每个 task 用一个 GPU，Slurm 不是”礼貌地请你用一块 GPU”，而是在操作系统层强制：

1. CUDA_VISIBLE_DEVICES：Slurm 在进程内设置该环境变量，让 CUDA 只看到列表里的设备（单节点有效）。
2. Linux cgroups 的 device allowlist：为 Task 0 建一个”沙箱”。

调试陷阱：若每个 GPU 上跑不同 task，看日志时每个错误都会指向 GPU 0（即使实际不是 GPU 0）， 因为每个进程看到的可见设备都从 0 重新编号。

2.2 启动与绑定步骤

Slurm 用 PMIx 编排启动后，应用侧初始化 PMIx 的步骤：

1. 连接 Slurm daemon
2. 初始化 PMIx，拿到 namespace 和全局 rank
3. 拿到 local rank
4. 用 local rank 选择具体的 GPU
5. 结束进程

• CUDA_VISIBLE_DEVICES 让进程看到过滤后、从 0 重编号的 GPU 集合。
• PMIX_LOCAL_RANK 是节点内 rank，是”把任务分配到该节点 GPU”的自然索引。
• 用 cudaGetDeviceCount 得知过滤后可见多少设备；需要原始物理索引时可解析 CUDA_VISIBLE_DEVICES。
• 深度调试用 cudaDeviceGetPCIBusId，把”可见设备 0”对应到 Slurm 实际分配的硬件 GPU。
• 绑定到本地 GPU 后启动 kernel；结束用 PMIx_Finalize(NULL, 0) 告诉 Slurm 完成。

3. NCCL 概述与初始化

3.1 什么是 NCCL

NVIDIA Collective Communications Library —— 面向 GPU 的库，让 GPU 间的集合通信又快又省心。

• 分布式训练中 GPU 持续交换张量（梯度/参数），NCCL 提供高度优化的原语： AllReduce、AllGather、ReduceScatter 等。
• 在 CUDA stream + 指针层面集成：你交给 NCCL 设备指针 + 一个 cudaStream_t， NCCL 自己在该 stream 上调度 GPU 工作（kernel + memcpy + 网络操作），与你的 kernel 通过普通 CUDA stream 顺序组合。
• 做 AllReduce 时，CPU 只是把一个 NCCL kernel 发到 GPU stream 上。

3.2 初始化：NCCL Unique ID

• 通信前，NCCL 需要一个”所有进程找到共同汇合点”的方式，它由 NCCL Unique ID 表示。
• 只有一个进程能生成该 ID；想加入组的每个进程都必须持有完全相同的 ID。
• 流程：Rank 0 让 NCCL 创建 Unique ID → 放进 PMIx 的 Key-Value Store（KVS）供他人读 → Rank 0 commit 推送出去 → fence 保证所有人都能读到 → 各 follower 查 KVS 并调 NCCL 初始化。

3.3 ncclGetUniqueId 与 ncclCommonInitRank

• ncclGetUniqueId：起始时各 GPU 不知周围其它 GPU 及如何经 NVLink/InfiniBand 到达它们。 为通信，NCCL 要建一个 communicator。ncclGetUniqueId 创建 ncclUniqueId 结构并记下第一块 GPU 的 IP， 之后每块 GPU 进来把自己的 IP 记进该结构。
• ncclCommonInitRank：设置阶段最贵、最复杂的函数——ID → 硬件连接：
  • 每个 rank 从 ncclUniqueId 提取 Rank 0 的 IP:Port；
  • 每个 rank 建一个标准 TCP socket 连 Rank 0；
  • Rank 0 等到收满 nranks 个连接——这是一个 barrier，Rank 799 慢则所有人等；
  • TCP 建立后不发数据，先发”自己的 rank、连接类型、设备”等元数据。

3.4 Rank 0 作为”架构师”：路由

• Rank 0 分析全局图，为 AllReduce 等操作找最高带宽、最低延迟路径：
  • 节点内优先用 NVSwitch；跨节点选特定 InfiniBand NIC；
  • 决定用 Ring（延迟优化）还是 Tree（带宽优化）。
• Rank 0 把”路由表”（发给谁、从谁收）发回每个 rank；节点内 rank 互相映射内存、 配置 NVSwitch 允许 GPU 直访；跨节点 rank 做 RDMA 握手，确保能不经 CPU 直接写对方内存缓冲。

3.5 ncclCommDestroy

• 拆除之前建立的硬件路径：把 communicator 标记为对后续 kernel 启动无效。
• 若 GPU 正在执行 NCCL kernel（如 stream 里的 AllReduce），不会立即抽走内存，靠内部引用计数。
• 释放：Init 时在 HBM 分配的 4–8MB scratchpad、CPU RAM 里用于 PCIe 的 staging 区、NVLink 映射与 InfiniBand Queue Pair。

4. NCCL 六大集合通信原语

在 H100 集群上，我们不按”发送/接收”想，而按跨 NVLink 结构的数据操作模式想。六大原语：

5. AI 的四种并行策略

5.1 Data Parallelism（数据并行）

• 最常见策略：模型在每个 GPU 上各有一份完整副本，全局 batch 切成 mini-batch，每 GPU 拿自己那份。
• 每 GPU 对各自数据做前向+反向；优化器更新前，必须保证所有副本一致 → 聚合（通常平均）各 GPU 的梯度。
• AllReduce 是数据并行最重要操作（每次反向结束发生）：GPU i 有梯度 gi，需每 GPU 都拿到平均梯度 (1/N)Σgi。
• Broadcast 用于训练开始或 checkpoint 恢复：Rank 0 初始化参数并广播给所有 rank，保证第 0 步所有副本数学上一致。

5.2 Tensor Parallelism（张量并行）

• 模型大到单 GPU 装不下时（训练时内存需求再 ×3–4：权重 + 梯度 + 优化器状态 + 激活），把矩阵切成块分到不同 GPU 算，算完再合起来。
• 切分方式：
  • 列切分（Column-wise）：用 Broadcast 把完整输入拷给各 worker，乘完用 AllGather 合并结果。
  • 行切分（Row-wise）：乘完用 AllReduce 求和得到最终结果。
  • 一个标准 Transformer block 的 TP 意味着 2 次 AllReduce（注意力子层 1 次 + FFN 子层 1 次）。

5.3 Sequence Parallelism（序列并行，SP）

• 标准 TP 只切重的矩阵乘（Linear 层），不切 Linear 之间的操作（LayerNorm、Dropout、GELU/SiLU）。
• 标准 TP 里，Linear 后 AllReduce 的输出是复制的——8 GPU 每块都存一份完整激活矩阵 [SeqLen, HiddenDim] 只为做 LayerNorm/Dropout。
• 但 LayerNorm 只作用于单个 token 的 Hidden 维，在 Sequence 维上独立——无需每块 GPU 复制完整序列，可把序列切分。
• 打破 AllReduce：数学上 AllReduce = ReduceScatter + AllGather。SP 把 LayerNorm/Dropout 注入通信循环： 先 ReduceScatter → 停 → 在分片数据上做 LayerNorm → 需要完整数据做下一次矩阵乘时才 AllGather。

5.4 Pipeline Parallelism（流水线并行）

• 若 TP 是”横向切一层”，PP 就是纵向切模型（切层堆）。
• 每 GPU 只存自己那几层的参数与优化器状态。
• 一次传大 batch 会大量 GPU 空等 → 把大 batch 切成微小 batch（micro-batch）：GPU 1 算完第一个 micro-batch 就传给 GPU 2、立刻算第二个。
• PP 只和流水线”邻居”通信 → 只用点对点 send/recv。

5.5 Expert Parallelism（专家并行，MoE）

• 与每个输入都用全部参数的”稠密”模型不同，MoE 只激活一小部分参数（专家）。
• 好处：参数量可 ×100（加更多专家），但每 token 只选 top-2 专家，算力开销仍较低。
• 专家分布在不同 GPU（GPU 1 持专家 A/B，GPU 2 持 C/D……）；token 需要别的 GPU 上的专家时经网络发过去 → 最终每 GPU 都要给其它每 GPU 发数据 = all-to-all 通信。
• 专用库（如 DeepEP，底层 NVSHMEM）擅长此场景：先跨节点高效打包发送，再用更快的 NVLink 在本地快速分发。

6. NCCL 各操作详解

6.1 NCCL 调用的通用格式

sendbuff, recvbuff, count, datatype  → 告诉 GPU 搬多少个元素、如何解释
op                                    → 执行什么数学操作
comm                                  → 通信句柄（存节点与集群的元数据）
stream                                → 指定异步行为

6.2 ncclBroadcast

• Root GPU 把一个张量发给 communicator 里其它所有 GPU。
• H100 DGX 上：Root GPU 把数据包发一次给 NVSwitch，NVSwitch 自己把它物理复制到其它 7 个端口（同时）。
• 因交换机负责复制，广播给 8 GPU 的时间 ≈ 发给 1 GPU 的时间。
• datatype 传 ncclFloat8e4m3 / ncclFloat8e5m2 时，NCCL 切换到 Hopper 专用 intrinsic 指令。

6.3 ncclReduce 与 Reduction

• Reduction：把每块 GPU 的数据用数学算子（sum/max）合并，结果存到某块（或所有）GPU。
• 旧系统要”水桶接力”式传数据求和；DGX H100 里 NVSwitch 自己做数学，把工作从 GPU 卸载。
• 通用 ncclReduce 通常用 NVLSTREE（NCCL 2.18+ 引入）处理”归约到单 root”的流（尤其跨多节点扩展时）。
• 用途：数据消费者集中（通常是 Rank 0）时用，如推理最后一层只有 Rank 0 需要完整 logits 做 argmax/top-k。
• count 是每 GPU 贡献的元素数，root 收到恰好 count 个元素。

6.4 ncclAllReduce

• 训练中最重要操作之一：N 块 GPU 合并 → 得全局结果 → 分发给 N 块 GPU。
• 8 块 GPU 同时把本地梯度从 HBM 推上 NVLink 通道、指向 NVSwitch → 交换机做 in-flight 归约 → 立即把结果多播回 8 块 GPU → 结果直接落到各 GPU 的 HBM3 目标缓冲。

ncclRedOp_t（归约算子）：

6.5 ncclReduceScatter

• 每 GPU 从完整梯度缓冲开始，跨 GPU 求和后再打散，每 GPU 只拿到最终和向量的一块。
• 输入：每 GPU 有 [A,B,C,D]；输出：GPU 0 持 Sum(A)、GPU 1 持 Sum(B)……
• recvcount 是输出缓冲的元素数（不是总输入），= 总元素 / GPU 数；Total Elements Reduced = recvcount × nranks。
• 不支持 “jagged” 数组：总向量 100、3 GPU 时 100/3 = 33.33 分不均——必须填充到 GPU 数的整数倍。

6.6 ncclAllGather

• 每 GPU 从不同片段开始，变换后每 GPU 拿到完整副本。
• 8 GPU 同时把梯度推进 NVSwitch 结构，NVSwitch 合并数据并路由到所有 GPU。
• 总总线利用率最大化，更接近理论 900 GB/s 聚合吞吐（不用等 ring 转一圈）。
• sendcount = 本 rank 贡献的元素数（本地切片大小）；每 rank 最终在 recvbuff 拿到所有 rank 的切片。

6.7 ncclAllToAll 与 MoE 问题

• 每 GPU 持有”要给其它每块 GPU”的数据并分别传输。
• 与常用复杂 Ring/Tree + SHARP 卸载的 AllReduce 不同，AllToAll 物理上更简单但带宽密集。
• 逻辑上是 N×(N−1) 次点对点（加自身），但作为一个集合操作协调优化。
• MoE 问题：NCCL 标准 AllToAll 是为静态、批量同步通信设计的，而 MoE 路由本质是动态、稀疏、延迟敏感的：
  • NCCL 是 CPU 驱动（Host API）、数据大小静态、要求连续块、且是阻塞的。
  • 专用库 DeepEP（底层 NVSHMEM） 更优：设备发起、细粒度数据搬移，能处理专家路由的不规则流量而不卡 GPU。

6.8 集合通信在 NVSwitch 里的优势小结

• Broadcast：交换机复制，8 GPU 广播 ≈ 1 GPU。
• Reduce/AllReduce：交换机做 in-flight 归约 + 多播回写。
• AllGather：交换机合并 + 路由，总线利用率最大化。

7. 点对点与分组：ncclSend/Recv 与 ncclGroup

7.1 ncclGroupStart / ncclGroupEnd

• 分布式系统常一个 CPU 线程管理多块 GPU；CPU 发 NCCL 指令时，若是阻塞调用会被卡住。 多数 NCCL 集合调用（如 ncclAllReduce）对 CPU 是异步的，更阻塞的往往是 communicator 初始化（ncclCommInitRank）。
• 若一个 host 线程逐条给多 GPU 发 NCCL 调用，可能造成部分参与方已开始、其它还没发出的中间态 → 死锁。
• ncclGroupStart()/ncclGroupEnd() 让你批量一组 NCCL 调用：组内收集、ncclGroupEnd() 时一起提交，避免部分启动问题。

7.2 ncclSend / ncclRecv

• ncclSend：非阻塞地把数据从某 GPU 发到另一 GPU，几乎总与目标设备的 ncclRecv 配对。
  • peer：目标 GPU 的 rank。
  • NCCL 自动找两 GPU 间最快物理路径（NVLink / PCIe / IB RDMA）。
• ncclRecv：告诉某 GPU 分配空间、等待指定 peer 的数据。异步（入队后 CPU 立刻恢复）。
  • 一个 ncclRecv 必须有匹配的 ncclSend；多个 P2P 传输须在各 GPU 上按兼容顺序入队，避免循环依赖。

8. 总结与学习衔接

8.1 核心脉络速记

8.2 全课程收束

本文档是 H100 系列（Lesson 1–10）的最后一课，把前 9 课的成果串成完整图景：

单卡 H100（TMA / Tensor Core / WGMMA 异步化）  →  L1–L7
  → 高性能单 kernel 设计（warp 特化 / 流水线 / 调度 / epilogue）  →  L8
  → 多 kernel 交接（依赖启动）  →  L8.2
  → 多 GPU 硬件互联（NVLink/NVSwitch/Rail 网络）  →  L9
  → 多 GPU 编程（Slurm/PMIx/NCCL/四种并行）  →  L10

8.3 一句话记忆

多 GPU 训练 = Slurm+PMIx 把进程铺满集群 → NCCL 用 NVSwitch 做 in-flight 集合通信 → 按模型形态选并行策略（数据/张量/流水线/专家）→ 把 AllReduce 拆成 ReduceScatter+AllGather 并在中间融合算子 → 用 NCCL 的六大原语在 900 GB/s 的 mesh 上最大化吞吐。

参考来源：10. Multi GPU Part 2.pdf（Lesson 10，60 页）。