通信与计算重叠：NCCL collective 与 kernel overlap

单卡算子之外，大模型必然多卡。多卡引入通信——梯度同步、激活传递、专家路由——而通信带宽远低于显存带宽，过 NVLink 或 PCIe/网络。如果通信和计算串行执行，通信时间就是纯开销，GPU 在等数据时空转。重叠（overlap） 是把通信藏到计算后面、让两者并发的核心手段。这一篇从 kernel 内到分布式，讲三个层次的重叠。它与 大模型基础设施系列 的并行篇互补——那里讲并行策略，这里讲算子和执行层面怎么实现重叠。

一、重叠的三个层次

重叠这个思想在 GPU 上贯穿多个尺度，机制不同但目标一致——让等待的时间被有用的工作填满：

flowchart TD
  A["层次一：kernel 内<br/>cp.async 异步加载<br/>访存与计算重叠"] --> B["层次二：kernel 间<br/>CUDA stream 并发<br/>独立 kernel / 拷贝与计算重叠"]
  B --> C["层次三：分布式<br/>NCCL collective 与计算重叠<br/>通信藏到反向/前向后面"]

这其实是 occupancy 篇 延迟隐藏思想的放大：warp 级用多 warp 隐藏访存延迟，kernel/分布式级用 stream 和异步通信隐藏更大尺度的延迟。

二、层次一：kernel 内的 cp.async

Ampere（本系列测试卡所属架构）引入 cp.async 指令，让数据从 global memory 异步拷贝到 shared memory，不经过寄存器、不阻塞后续计算。这是 GEMM/Tensor Core kernel 实现”算当前 tile 时预取下一个 tile”（double buffering）的硬件基础。

// 概念：发起异步拷贝，立即返回；之后用 cp.async.wait_group 等它完成
cp.async.cg.shared.global [shared_ptr], [global_ptr], 16;
cp.async.commit_group;
// ... 这里可以做当前 tile 的 MMA 计算 ...
cp.async.wait_group 0;   // 需要数据时再等

没有 cp.async 时，加载 tile 要先把数据读进寄存器再写进 shared，占寄存器、且加载和计算难以重叠。有了它，GEMM 篇 提到的”通往 cuBLAS”优化里的预取才能高效实现——这也是 CUTLASS 的 mainloop 大量使用 cp.async 做软件流水的原因。回到 Tensor Core 篇 的结论：Tensor Core 算得太快，必须用异步加载把它喂饱，cp.async 就是喂数据的关键工具。

三、层次二：CUDA stream

stream 是 GPU 上的命令队列。同一个 stream 内的操作顺序执行；不同 stream 的操作可以并发（只要硬件资源够）。这让独立的工作互相重叠：

cudaStream_t s1, s2;
cudaStreamCreate(&s1); cudaStreamCreate(&s2);
cudaMemcpyAsync(d_a, h_a, n, cudaMemcpyHostToDevice, s1);  // 拷贝在 s1
kernel<<<g, b, 0, s2>>>(d_x);                               // 计算在 s2，与拷贝并发

典型用途：

• 拷贝与计算重叠：把 H2D/D2H 拷贝放一个 stream、计算放另一个，PCIe 传输和 GPU 计算同时进行（需要 pinned host memory 才能真正异步）。
• 独立 kernel 并发：多个小 kernel 各占不满 GPU 时，放不同 stream 让它们并发填满 SM。
• 流水线：把一个大任务切成块，块的拷贝、计算、回传在不同 stream 上流水。

stream 之间的依赖用 event 表达（cudaStreamWaitEvent）。Nsight Systems（第 08 篇）的时间线就是用来看 stream 有没有真正并发、有没有意外的同步点把并发打断。

四、层次三：分布式里的通信-计算重叠

多卡训练的通信主要是 NCCL（NVIDIA Collective Communications Library） 提供的集合通信：all-reduce（梯度求和）、all-gather、reduce-scatter、all-to-all（MoE 路由）。这些 collective 本身是高度优化的 kernel（如 ring all-reduce 把数据切环传递），但关键在于和计算重叠。

数据并行的梯度 all-reduce 是经典例子。反向传播按层从后往前算梯度，而某一层的梯度一算完，就可以立即开始它的 all-reduce——不必等所有层都算完。于是：

反向计算:  [layer N] [layer N-1] [layer N-2] ...
all-reduce:          [grad N]    [grad N-1]   ...   ← 与后续层的反向重叠

后面层的反向计算和前面层梯度的通信并发进行，通信时间被反向计算盖住。这就是 PyTorch DDP 的 gradient bucketing + 异步 all-reduce 在做的事。其他例子：

• 张量并行：层内的 all-reduce/all-gather 尽量和计算重叠，但层内重叠空间小、对互联带宽（NVLink）要求高。
• 流水线并行：micro-batch 之间流水，用气泡填充隐藏跨 stage 的通信。
• MoE：专家并行的 all-to-all 与专家计算重叠。

这些策略的完整讨论见 大模型基础设施系列的并行篇 和 互联篇；本篇强调的是它们都依赖底层的 stream/异步机制来实现重叠。

五、重叠的代价：资源争抢

重叠不是免费的。让通信和计算并发，意味着它们要共享 SM 资源——NCCL 的 collective kernel 也要占 SM 和带宽。几个现实约束：

• SM 争抢：通信 kernel 占用的 SM 越多，留给计算的越少。NCCL 可以配置占用的 SM 数（NCCL_MAX_NCHANNELS 等），需要在通信吞吐和计算资源间权衡。
• 带宽争抢：通信和计算可能都要访问 HBM，重叠时争抢显存带宽，未必能完全并行。
• 同步开销：过细粒度的重叠（太多小 collective）会被启动和同步开销吃掉收益，所以要 bucketing（把小梯度攒成大块再通信）。
• 依赖正确性：重叠引入并发，必须用 event/stream 依赖保证数据就绪，否则读到未完成的通信结果。

判断重叠效果还是靠 profiler 的时间线：理想情况下通信条目应该和计算条目在时间上重叠，而不是串在计算后面。

六、小结与下一步

• 多卡场景下通信不与计算重叠就是纯开销；重叠是把通信藏到计算后面的核心手段。
• 三个层次：kernel 内用 cp.async 让访存与计算重叠（喂饱 Tensor Core）、kernel 间用 stream 让拷贝/独立 kernel 并发、分布式用 NCCL collective 与反向计算重叠。
• 数据并行的梯度 all-reduce 与后续层反向重叠是经典模式，依赖 bucketing 和异步通信。
• 重叠有代价：SM 和带宽争抢、同步开销，需要权衡和 profiler 验证。

算子实现的技术面到此覆盖完。最后三篇转向工程化：先讲 算子库工程：dispatch、autotune cache 与 JIT，把单个 kernel 变成可维护的库。