0. 总览两篇论文的定位差异维度NCCL M2N2610.07516T-CCL2610.07098层级语义层—— 新增 集合原语layout-aware reshard collective引擎层—— 替换 通信执行引擎TMA 卸载解决问题跨 group 的 M-to-N 重排没有原生集合表达集合通信占用过多 SM 资源挤压并发计算作者NVIDIA 官方NCCL/NCCL Extensions 团队Chalmers/Linnaeus学术团队含 10-07 日报 MoE 论文作者 Minyu Cui形态作为nccl_m2n开源于 NVIDIA/nccl-extensions开源于 github.com/ooverlap/ooverlapSC26 研讨会一句话M2N 回答 该传哪些字节、走哪条路T-CCL 回答 用谁去搬字节、谁去做归约—— 两者可叠加。1. NCCL M2N1.1 问题背景Tensor 使用不同并行布局训练侧通常是 TPPPEP生成侧是 TPDP每步训练后权重需从训练布局重排reshard到生成布局 —— 即M-to-N 布局重分布M 个源 rank 持有一份 layoutN 个目的 rank 需要另一份。现有 NCCL 没有任何集合能表达 两个任意 layout 端点 不相交 process group 的转移。两种 workaround 都昂贵flat direct sends每个源贡献发给所有目的副本 → 网络流量 × 副本数gather-then-broadcast汇聚到 root 再广播 → root 注入成为瓶颈且发送了目的端不需要的完整张量。量化代价256 GPU DeepSeek-V3legacy all-gatherbroadcast 的weight sync 占 RL step time 的 29.4%。1.2 创新点First-class M-to-N reshard 集合原语以 (mesh, placement) 描述符为输入调用时自动推导传输计划 —— 不再需要手写 per-pair 调度。Layout-derived traffic只传目的 layout 真正需要的字节区间相交计算消除 over-delivery。Hierarchical topology route每个目的 NVLink 域只注入一份副本域间 ring 转发、域内 NVLink 本地复制 —— 消除 replica-multiplied source egressleader 选择按负载均衡loads differ by at most one。Aggregate data-movement 模型T_SOL max(T_network, T_NVLink)给出速度上限可验证性。实测单 FFN-MoE 层传输7.9x9.8ms vs 77.3ms256-GPU DeepSeek-V3 NeMo-RLweight sync 5.78s→2.77s2.09xstep time -12.7%KL 误差基本持平0.001869 vs 0.001912。1.3 Kernel 侧主要 API论文公开了完整的 Python APIListing 1与 C 接口存在性device 侧复用 GIN/LSA 机制。整理如下表格层API说明Pythonm2n.Mesh([dims], start_rankbase)声明网格维度与 base rankmesh 内的 rank 是 communicator 局部连续区间Pythonm2n.Shard(d)/m2n.Replicate()placement 描述符Pythonm2n.reshard(src_t, dst_t, comm, stream, src_mesh, src_placements, dst_mesh, dst_placements)核心调用仅参与的端点传 buffer另一端点只传 layout 元数据CncclReshard(...)nccl_extensions 约定C 版本两端各传本地指针 mesh/placement 描述符DevicencclGetP2pPtr、GINput() signal已有数据面域内 LSA 直读、跨域 GIN 一侧 RDMA设计约束当前实现要求 M_s 与 M_d 在 communicator 内不相交每 layout 至多一个 sharded mesh axis协议错误为 fail-stop。1.4 代码样例示例 ATP 同维合并 目的端复制TP4 → TP2, DP2import nccl.m2n as m2n import torch # 训练器TP4, DP14 个 rank0-3每 rank 一行 [1:2] 等 src_mesh m2n.Mesh([1, 4], start_rank0) # 生成器TP2, DP24 个 rank4-7rank4/6 持 [0:2]rank5/7 持 [2:4] dst_mesh m2n.Mesh([2, 2], start_rank4) # 每个 rank 在自己的 role 里提供 buffer不参与的一端传 None src_tensor trainer_weight if is_trainer else None dst_tensor generator_weight if is_generator else None m2n.reshard( src_tensor, dst_tensor, comm, streamtorch.cuda.current_stream(), src_meshsrc_mesh, src_placements[m2n.Replicate(), m2n.Shard(0)], # DP1 轴复制、TP 轴 shard 维度0 dst_meshdst_mesh, dst_placements[m2n.Replicate(), m2n.Shard(0)], # DP2 复制、TP2 shard 维度0 ) # 完成后G4/G6 持有 [0:2,0:8]G5/G7 持有 [2:4,0:8]DP 副本由 NVLink 域内复制产生示例 BPP 逐阶段分解TP2,PP2 → TP1,PP1—— 论文原样# rank 4 同时加入两个 PP communicator作为 c-rank 2 for stage, pp_comm in local_stage_comms: src_mesh m2n.Mesh([2], start_rank0) # 训练器 TP 轴本 PP stage 的 rank dst_mesh m2n.Mesh([1], start_rank2) # 生成器本 PP comm 的 c-rank 2 src_tensor local_stage_weights[stage] if is_trainer else None dst_tensor generator_weights[stage] if is_generator else None m2n.reshard( src_tensor, dst_tensor, pp_comm, streamstream, src_meshsrc_mesh, src_placements[m2n.Shard(0)], dst_meshdst_mesh, dst_placements[m2n.Replicate()], )Device 侧数据面示意—— 数据面基于 GIN跨域注入可表达为// 示意leader 注入 域内复制。真实实现见 nccl-extensions 的 nccl_m2n 组件。 // 每个源贡献只注入到其目的 shard 的 leader负载均衡选定 // 域内其余副本由 NVLink 复制LSA peer load/store 或 TMA bulk 复制。 if (is_leader_for_my_contrib) { // GIN: GPU-initiated one-sided RDMA put signal ncclGinPut(src_buf contrib_off, dst_leader_buf dst_off, bytes, stream); ncclGinSignal(dst_leader, stream); } __syncthreads(); // 域内 fan-outNVLinkleader → 同域其余 ρ-1 个副本 for (int r 1; r rho; r) { peer_copy_via_nvlink(dst_leader_buf, dst_replica_buf, bytes); // LSA / TMA }1.5未来方向on-the-fly quantization权重传输前量化融合进 staging kernel省一次中间 buffer—— 这与压缩式集合通信ZipCCL互补elastic communicatorncclCommGrow/Shrink让 reshard 计划可原地更新 —— 弹性 RL 训练。2. T-CCLTMA 卸载的机内集合通信2.1 问题背景NCCL 的机内集合 kernel 依赖大量 GPU 线程做数据搬移与归约 →SM 资源占用大在 compute-communication overlapTP 层中 GEMM 与集合并发场景通信 kernel 挤占 SM拖慢 GEMM。线程驱动的复制可灵活归约但吃 SMcopy engine 异步但不支持元素级归约TMA 两者兼具异步 bulk 搬移 异步 bulk reductionPTX 定义 add/min/max 元素级原子归约单线程即可发起。2.2 创新点TMA 全卸载数据移动 归约都交给 TMA而非仅搬移每个集合执行为流水化的异步 TMA 操作序列。Reduction 本地化不做跨设备原子归约 ——shard-in 阶段 owner 从 peerload数据并本地归约跨设备归约比 load 后本地归约贵得多同步点不随消息规模增长。Shard-out 的 load once, reuse for all peer storesowner 把最终 shard 一次 load 进 SMEMTMA 复用同一份 SMEM 数据 store 到所有 peer—— 单个 bulk load 服务全部目的。兼容用户 buffer对比 NCCL symmetric kernels 需要VMM 分配的 bufferT-CCL 接受 PyTorch 等框架管理 buffer—— 这降低了接入成本。实测unrestricted 下最高2.4xvs NCCL、restricted CTA8/9 CTA下最高3.42x与 symmetric NCCL 相当或更优GEMM 重叠案例 1.12→1.25x2 GPU、1.04→1.14x4 GPUvLLM 端到端最高1.31xdecode-heavy 收益最大。2.3 实现方案编程模型Grouprank支持 AllReduce/ReduceScatter/AllGather单进程一个进程建全组与多进程每进程建本地视图IPC 采用 ThunderKittens utilities两种模式in-place 优先省任务、省多进程下的 buffer 信息交换接口语义对齐 NCCL便于替换。Task-based planning逻辑张量分 P 个 shard 分配给 P 个 rankplanner 只在最低 ID 的 rank 上运行生成逻辑任务op、执行 rank、源 / 目的区间、依赖经预分配 buffer 广播给其他 rank每个 rank 独立把逻辑任务lower为物理任务解析地址与传输机制—— 顶层拓扑规划与设备级执行解耦。TMA engine物理任务 一系列 windowchunk 拆解进 SMEM 多级流水线论文图 24 stage 处理 6 chunkscp.async.bulk.tensor系列指令 mbarrier 完成信号launch 配置可 sweep 调优test/sweep_tma_collectives.py。2.4 Kernel 侧主要 APIT-CCL 的对外接口是 NCCL 语义的 host APIkernel 侧的核心是 PTX TMA 原语组合。整理表格层API / 原语说明Hosttcl::Group::create(...)/rank()建组单 / 多进程模式IPC 用 ThunderKittens 工具Hosttcl::allreduce(sendbuf, recvbuf, count, op, stream)等NCCL 语义支持 in/out-of-placeHosttcl::reducescatter(...)/tcl::allgather(...)三种集合Device (PTX)cp.async.bulk.tensor.2d.global.shared::cta.global.mbarrier::complete_tx::bytesTMA bulk 搬移global↔SMEM含 peer memoryDevice (PTX)cp.async.bulk.tensor.2d.global.shared::cta.global.mbarrier::complete_tx::bytesreduce 变体如.red.op.addTMA 异步归约add/min/max单线程发起、硬件执行Device (PTX)mbarrier.arrive.expect_tx/mbarrier.try_wait.parity完成信号 / 流水线阶段同步2.5 代码样例Host 侧NCCL 语义示意#include tcl/tcl.hpp // 单进程模式一个进程创建包含全部 GPU 的组 tcl::Group group tcl::Group::create({0, 1, 2, 3}, /* single-process */ true); int rank group.rank(); // PyTorch 管理的普通 buffer 即可对比 NCCL symmetric 需 VMM 分配 auto in torch::zeros({4096, 4096}, torch::dtype(torch::kFloat16).device(rank)); auto out torch::empty_like(in); // in-place AllReduceT-CCL 内部为每个 shard 生成 shard-in(load本地归约) // 与 shard-out(load-once 对 peer 复用 store) 任务交给 TMA engine 流水化执行 group.allreduce(in.data_ptr(), out.data_ptr(), in.numel(), tcl::ReductionOp::SUM, /* stream */ at::cuda::getCurrentCUDAStream());Device 侧 TMA 引擎示意chunk 流水线核心// 示意shard-in 阶段——owner 从 peer load 贡献并本地归约 // 每个 chunkTMA bulk load (peer global - SMEM) TMA bulk reduce (SMEM - owner global) __global__ void tma_shard_in_kernel(...) { // 1. 发起异步 TMA load单线程即可目标 SMEM stage[i] cp_async_bulk_tensor_2d(peer_gmem_desc, smem_desc[stage], mbar[stage], bytes); // 2. 等待 load 完成 mbarrier_wait(mbar[stage]); // 3. 发起 TMA 异步归约SMEM 分片 - owner 的 global shard硬件执行 add cp_async_bulk_tensor_reduce_add(owner_gmem_desc, smem_desc[stage], mbar[stage1], bytes); // 4. 流水线推进stage 复用双/多缓冲 mbarrier parity 翻转 } // 示意shard-out 阶段——load once, reuse for all peer stores __global__ void tma_shard_out_kernel(...) { // owner 把最终 shard 一次 load 进 SMEM cp_async_bulk_tensor_2d(owner_gmem_desc, smem_desc, mbar_in, bytes); mbarrier_wait(mbar_in); // 同一份 SMEM 数据对每个 peer 的对应区域发起异步 store复用不重复 load for (int p 0; p num_peers; p) { cp_async_bulk_tensor_2d(peer_gmem_desc[p], smem_desc, mbar_out[p], bytes); } // 完成信号仅用于集合边界同步不随消息规模增长 }3. 交叉对比与组合建议M2N × T-CCL 可叠加M2N 决定 每个源贡献进入哪个目的 leader、域内如何复制T-CCL 可做 域内复制用什么引擎搬——M2N 的 NVLink fan-out 若用 TMA bulk 复制 本地归约就是两者结合。NCCL 2.27 symmetric kernels 已在 Blackwell 默认启用 TMA 路径方向一致但受 VMM buffer 限制。与 mKernel/PK/ZipCCL 的关系mKernelSM 角色化compute/communication 分工 tile 级传输 —— 与 T-CCL少占 SM 同一目标一个在 kernel 内部做分工、一个用硬件引擎替换ParallelKittens多播内存 / 网内归约 ——T-CCL 论文明确对比PK 需要 NVSwitch 机制T-CCL 只需 GPU 内已存在的 TMAZipCCL压缩减流量 —— 与 M2N 的 on-the-fly quantization 未来工作同族传输前降体积。研究最值得借鉴的三点M2N 的 区间相交 → 确定性全局计划是 layout-aware 通信的标准建模法DTensor/GSPMD 之上薄薄一层即可实现无需改框架T-CCL 的 reduction 本地化 load-once-store-many是 少占用 SM 的通用原则可直接套用到 MoE 的 all-to-all 分段传输结合 10-07 的 MoE tile 级论文SM 分区里通信 kernel 用 TMA 而非线程驱动两篇的数据移动 SOL 模型T_SOL max(网络, NVLink)为你的 benchmark 提供了可验证上界方法论。