高性能计算通信深度学习分布式训练【免费下载链接】ncclOptimized primitives for collective multi-GPU communication项目地址https://gitcode.com/gh_mirrors/nc/nccl点击查看免费下载导读本指南以 NCCL 开源仓库 docs/examples/05_symmetric_memory 目录下的两个最小示例为核心系统讲解 NCCL 对称内存Symmetric Memory窗口的完整生命周期如何用ncclMemAlloc分配 VMM 兼容缓冲区、如何用ncclCommWindowRegister以NCCL_WIN_COLL_SYMMETRIC标志注册对称窗口以及如何结合CTAPolicy2Copy Engine实现零 SM 占用的 AllGather。读完本文你将掌握对称内存的 C/Python 双语言实现、正确的资源释放顺序、常见故障排查以及从源码层面理解其底层运行原理。背景为什么需要对称内存窗口当所有 rank 使用一致的缓冲区布局相同大小、由 NCCL 分配器分配时NCCL 可以应用经过优化的通信算法降低延迟并提升带宽。对称内存窗口Symmetric Memory Windows正是为此设计的机制每个 rank 向通信域注册自己的窗口NCCL 在集体通信时即可利用这种对称性。从源码结构看这一机制建立在 CUDA Virtual Memory ManagementVMMAPI 之上。ncclMemAlloc的实现src/allocator.cc#L16-L107完整展示了这一流程通过cuMemCreate分配物理内存并优先尝试CU_MEM_HANDLE_TYPE_FABRIC句柄CUDA_VERSION 12030且设备支持时否则回退到CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR通过cuMemAddressReserve预留虚拟地址范围通过cuMemMap将虚拟地址映射到物理分配通过cuMemSetAccess为所有支持 P2P 的设备授予读写权限若 CUDA 版本低于 11.3 或上述流程失败则回退到普通的cudaMalloc。正是这种 VMM 分配方式使得缓冲区具备跨设备访问能力成为对称窗口注册的前提。示例总览与构建方式仓库中该目录包含两个示例顶层 README示例模式核心 API关键特性01_allreduce每 rank 注册对称窗口并用于集体通信ncclCommWindowRegister、ncclCommWindowDeregister、ncclMemAlloc、ncclAllReduce大规模集体通信、一致缓冲区布局、正确注销与清理02_allgather对称窗口 Copy EnginencclCommInitRankConfig、ncclCommWindowRegister、ncclCommWindowDeregister、ncclMemAlloc、ncclAllGatherCTAPolicy2启用 Copy Engine、集体操作零 SM 占用、计算与通信真正重叠、大消息峰值带宽更高小消息延迟略高快速构建在 docs/examples/05_symmetric_memory 目录下按示例目录名构建# 按目录名构建示例 make 01_allreduce make 02_allgather该目录的 Makefile 还提供all、clean、test、install、help等目标test会先构建全部示例再逐个执行。若想构建整个示例集可在 docs/examples 顶层执行make会依次构建 01_communicators 到 07_kernel_fusion 共七个分类。单独构建与运行 C 示例# 构建并运行对称内存 AllReduce cd 01_allreduce/c make ./allreduce_sm # 构建并运行对称内存 Copy Engine AllGather cd 02_allgather/c make ./allgather_ceC 变体的构建命令支持以下可选参数见 01_allreduce/c/README.md 与 02_allgather/c/README.mdmake [MPI1] [MPI_HOMEpath-to-mpi] [NCCL_HOMEpath-to-nccl] [CUDA_HOMEpath-to-cuda]默认按 pthread 模式编译运行方式为[NTHREADSN] ./allreduce_smNTHREADS控制线程数默认取自 GPU 数量参见 docs/examples/common/src/utils.cc传入MPI1编译为 MPI 模式后用mpirun -np num_processes ./allreduce_sm运行适用于多节点场景。Python 变体每个示例目录下的python/子目录提供 nccl4py mpi4py 实现其运行方式与 README 说明一致详见 01_allreduce/python/README.md 与 02_allgather/python/README.mdpython -m venv .venv source .venv/bin/activate pip install -r requirements-cu12.txt # 或: requirements-cu13.txt mpirun -np num_processes python allreduce_sm.py # AllReduce 示例 mpirun -np num_processes python allgather_ce.py # AllGather Copy Engine 示例运行时环境要求C 变体可用 pthreads 或 MPI 构建多节点运行需要 MPIPython 变体为单进程启动每个 MPI 进程一张 GPU。对称内存的收益在单节点上最为自然多 GPU 时效果最佳但单张可见 GPU 也能运行。Copy Engine 的收益仅适用于单个 (MN)NVL 域内部CTAPolicy2在跨域场景下不被支持。示例一AllReduce with Symmetric Memory Windows代码走读C 实现完整源码位于 docs/examples/05_symmetric_memory/01_allreduce/c/main.cc其生命周期分为 7 步Step 1初始化 NCCL CommunicatorncclUniqueId nccl_unique_id; if (my_rank 0) { NCCLCHECK(ncclGetUniqueId(nccl_unique_id)); } util_broadcast(0, my_rank, nccl_unique_id); // 广播给所有 rank CUDACHECK(cudaSetDevice(local_device)); ncclComm_t comm; NCCLCHECK(ncclCommInitRank(comm, total_ranks, nccl_unique_id, my_rank));Step 2用 NCCL 分配器分配缓冲区对称窗口的硬性前提size_t count 1024 * 1024; // 1M 个元素 size_t size_bytes count * sizeof(float); void *d_sendbuff, *d_recvbuff; NCCLCHECK(ncclMemAlloc(d_sendbuff, size_bytes)); NCCLCHECK(ncclMemAlloc(d_recvbuff, size_bytes));Step 3注册对称内存窗口集体调用所有 rank 必须参与ncclWindow_t send_win, recv_win; NCCLCHECK(ncclCommWindowRegister(comm, d_sendbuff, size_bytes, send_win, NCCL_WIN_COLL_SYMMETRIC)); NCCLCHECK(ncclCommWindowRegister(comm, d_recvbuff, size_bytes, recv_win, NCCL_WIN_COLL_SYMMETRIC));Step 4-5初始化数据并执行 AllReduce// 每个 rank 用自身 rank 值填充缓冲区便于后续验证 for (size_t i 0; i count; i) h_data[i] (float)my_rank; CUDACHECK(cudaMemcpy(d_sendbuff, h_data, size_bytes, cudaMemcpyHostToDevice)); NCCLCHECK(ncclAllReduce(d_sendbuff, d_recvbuff, count, ncclFloat, ncclSum, comm, stream));Step 6验证结果每元素期望值为所有 rank 之和total_ranks * (total_ranks - 1) / 2用cudaMemcpy拷回主机后逐元素比对容差 0.001。Step 7按正确顺序清理顺序至关重要free(h_data); free(h_result); // 1. 释放主机内存 NCCLCHECK(ncclCommWindowDeregister(comm, send_win)); // 2. 注销窗口 NCCLCHECK(ncclCommWindowDeregister(comm, recv_win)); NCCLCHECK(ncclMemFree(d_sendbuff)); // 3. 释放设备内存 NCCLCHECK(ncclMemFree(d_recvbuff)); NCCLCHECK(ncclCommFinalize(comm)); // 4. 终结并销毁 communicator NCCLCHECK(ncclCommDestroy(comm)); CUDACHECK(cudaStreamDestroy(stream)); // 5. 销毁 CUDA 流4 GPU 下的期望输出Starting AllReduce example with 4 ranks Rank 0 communicator initialized using device 0 ... Symmetric Memory allocation Rank 0 allocating 4.00 MB per buffer ... Starting AllReduce with 1048576 elements (4 MB) AllReduce completed successfully Verification - Expected: 6.0, Got: 6.0 Results verified correctly Rank 0 symmetric memory windows deregistered ... All resources cleaned up successfully Example completed - demonstrated symmetric memory lifecycle性能收益与适用场景直接 P2P 访问GPU 可直接访问对端内存降低延迟消除中间拷贝与协议开销LSALoad-Store Accessible内存可通过直接 load/store 操作访问为 NCCL Device API 的 kernel 级通信提供基础。适用场景单节点内 GPU 可互相直连、对延迟敏感、在相同缓冲区上反复执行集体操作、使用 NCCL Device API 进行 kernel 级通信。示例二AllGather with Symmetric Memory and Copy Engine启用 Copy EngineCTAPolicy202_allgather/c/main.cc 与示例一的差异在于初始化阶段// 配置 NCCL 使用 Copy EngineCTAPolicy2实现零 SM 占用 // 等效方式设置环境变量 NCCL_CTA_POLICYZERO ncclConfig_t config NCCL_CONFIG_INITIALIZER; config.CTAPolicy 2; ncclComm_t comm; NCCLCHECK(ncclCommInitRankConfig(comm, total_ranks, nccl_unique_id, my_rank, config));从源码看CTAPolicy的取值定义在 src/nccl.h.in#L70-L73#define NCCL_CTA_POLICY_DEFAULT 0x00 #define NCCL_CTA_POLICY_EFFICIENCY 0x01 #define NCCL_CTA_POLICY_ZERO 0x02即示例中的CTAPolicy2对应NCCL_CTA_POLICY_ZERO。环境变量NCCL_CTA_POLICY的解析位于 src/init.cc#L73-L121支持DEFAULT、EFFICIENCY、ZERO三种取值以及传统数值 0/1/2。当策略为 ZERO 且 Copy Engine 可用ceAvailable || hierCeAvailable且缓冲区无系统内存段时NCCL 会在 src/enqueue.cc#L3054 走 Copy Engine 路径集体操作不再占用 SM。缓冲区与窗口注册接收缓冲区大小需为sendcount * total_ranks * sizeof(type)容纳所有 rank 的切片ncclMemAlloc(d_sendbuff, send_size_bytes); ncclMemAlloc(d_recvbuff, recv_size_bytes); ncclCommWindowRegister(comm, d_sendbuff, send_size_bytes, send_win, NCCL_WIN_COLL_SYMMETRIC); ncclCommWindowRegister(comm, d_recvbuff, recv_size_bytes, recv_win, NCCL_WIN_COLL_SYMMETRIC); NCCLCHECK(ncclAllGather(d_sendbuff, d_recvbuff, sendcount, ncclFloat, comm, stream));期望输出要点4 GPU 场景下每个 rank 分配 4.00 MB 发送缓冲与 16.00 MB 接收缓冲验证阶段依次检查 Segment 03期望值分别为 0.0、1.0、2.0、3.0即接收缓冲的每个分段都持有对应 rank 贡献的值。Copy Engine 的优势与适用场景零 SM 占用单个 (MN)NVL 域内集体操作使用 Copy Engine 而非 SM释放计算资源计算通信重叠实现通信与 GPU 计算 kernel 的真实重叠更高峰值带宽大消息尺寸下峰值带宽更高小消息延迟略高。最适用场景需要重叠通信与计算的应用程序、SM 资源紧张的场景、无算术操作的大规模集体通信AllGather、AlltoAll、Gather、Scatter。Python 变体nccl4py mpi4py对称内存 AllReduceallreduce_sm.py每进程一张 GPU通过MPI.COMM_TYPE_SHARED拆分节点内 rank 得到local_rank并以cp.cuda.Device(local_rank).use()设为当前设备Unique ID 分发rank 0 调用nccl.get_unique_id()用comm_world.bcast()广播对应 C 的ncclGetUniqueId 广播Communicator 初始化nccl.Communicator.init(nranksmpi_size, rankmpi_rank, unique_id...)对应ncclCommInitRankNCCL 托管 CuPy 数组nccl.cupy.empty(...)创建由 NCCL 内存管理、可直接注册为对称窗口的数组窗口注册comm.register_window(buf, flagsnccl.WindowFlag.CollSymmetric)集体调用对应ncclCommWindowRegisterNCCL_WIN_COLL_SYMMETRICAllReducecomm.allreduce(send_buf, recv_buf, nccl.SUM)注销comm.destroy()自动注销所有已注册窗口对应ncclCommWindowDeregister。对称内存 AllGather Copy Engineallgather_ce.py差异仅在配置nccl.NCCLConfig(cta_policynccl.CTAPolicy.ZERO)选择零 SM 策略对应 C 的config.CTAPolicy 2并通过Communicator.init(..., configconfig)传入对应ncclCommInitRankConfig接收缓冲大小为sendcount * mpi_sizecomm.allgather(send_buf, recv_buf)在对称内存 CTAPolicy.ZERO下运行于 Copy Engine。如何选择正确的模式场景具有一致内存模式的大规模训练解决目标在受支持的平台上实现低延迟、高带宽的集体通信依赖pthread 或 MPI。对称内存要求内存通过 CUDA VMM API 分配并注册到 NCCL。是否启用 Copy Engine取决于你是否需要把 SM 腾出来跑计算 kernel需要通信计算重叠时选择02_allgather的CTAPolicy2方案否则01_allreduce的默认策略即可。源码级原理窗口注册与对称性检查从 src/dev_runtime.cc#L1331-L1373 可以看到ncclCommWindowRegister的实现要点校验buff/size/win参数合法性若 communicator 不支持对称内存!comm-symmetricSupport直接返回成功窗口为空将注册任务封装为ncclDevrRegTask并入队regTaskQueue并通过ncclGroupCommJoin加入组操作——这也是它作为集体调用的实现依据。窗口标志定义于 src/nccl.h.in#L63-L68#define NCCL_WIN_DEFAULT 0x00 #define NCCL_WIN_COLL_SYMMETRIC 0x01 #define NCCL_WIN_STRICT_ORDERING 0x02 #define NCCL_WIN_REQUIRED_ALIGNMENT 4096NCCL_WIN_COLL_SYMMETRIC要求所有 rank 在集体操作中一致地提供对称缓冲区窗口对齐要求为 4096 字节。常见问题与解决方案窗口注册失败缓冲区必须来自 VMM 兼容分配器如ncclMemAlloc()不能使用cudaMalloc()分配错误检查 NCCL 版本需 2.27与可用内存注销顺序必须先注销窗口再释放内存或销毁 communicator对称性要求所有 rank 必须一致地使用NCCL_WIN_COLL_SYMMETRIC/WindowFlag.CollSymmetric内存泄漏用ncclMemAlloc分配的缓冲区必须用ncclMemFree释放Copy Engine 不受支持CTAPolicy2仅在单个 (MN)NVL 域内生效。延伸阅读对称内存 AllReduce 的 C 源码01_allreduce/c/main.cc运行说明见 01_allreduce/c/README.mdCopy Engine AllGather 的 C 源码02_allgather/c/main.cc运行说明见 02_allgather/c/README.mdNCCL 设备 API、用户缓冲注册等更深入的用法可参考仓库 docs/examples 的 04_user_buffer_registration、06_device_api 等示例分类用户指南位于 docs/userguide。赞分享高性能计算通信深度学习分布式训练【免费下载链接】ncclOptimized primitives for collective multi-GPU communication项目地址https://gitcode.com/gh_mirrors/nc/nccl点击查看免费下载相关推荐NCCL 对称内存 AllReduce 实战用 nccl4py mpi4py 跑通 Symmetric Memory 窗口通信NCCL 对称内存 AllReduce 实战用 nccl4py mpi4py 跑通 Symmetric Memory 窗口通信 导读 本文基于 NCCL高性能计算通信深度学习分布式训练NCCL 对称内存Symmetric MemoryAllReduce 实战指南从窗口注册到优化通信NCCL 对称内存Symmetric MemoryAllReduce 实战指南从窗口注册到优化通信 对称内存Symmetric Memory是 NCC高性能计算通信深度学习分布式训练NCCL C 语言对称内存Symmetric MemoryAllReduce 实战从窗口注册到性能原理NCCL C 语言对称内存Symmetric MemoryAllReduce 实战从窗口注册到性能原理 导读 本文基于 NCCL 官方示例 docs/ex高性能计算通信深度学习分布式训练上一篇antd Select 大数据量场景下的虚拟滚动原理、参数配置与关闭时机下一篇CANN ops-nn 算子解析HardSwishGrad 反向梯度算子的原理、约束与 aclnn 调用实践创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
